diff --git a/cuda_bindings/cuda/bindings/_lib/param_packer.h b/cuda_bindings/cuda/bindings/_lib/param_packer.h index 160ef5f7c92..5cbf6544fb4 100644 --- a/cuda_bindings/cuda/bindings/_lib/param_packer.h +++ b/cuda_bindings/cuda/bindings/_lib/param_packer.h @@ -7,6 +7,8 @@ #include #include #include +#include +#include static PyObject* ctypes_module = nullptr; @@ -69,7 +71,24 @@ static void populate_feeders(PyTypeObject* target_t, PyTypeObject* source_t) { m_feeders[{target_t,source_t}] = [](void* ptr, PyObject* value) -> int { - *((int*)ptr) = (int)PyLong_AsLong(value); + // Bound to the 32-bit int slot, NOT to `long`. AsLongAndOverflow's + // `overflow` only flags values outside `long`, which is 64-bit on + // LP64 (Linux/macOS) -- so a value in 2**31..2**63 comes back with + // overflow==0 and would be silently truncated by (int)v. The + // explicit INT_MIN/INT_MAX check rejects it. When overflow!=0, v is + // the -1 sentinel (not the real value), so that case must be caught + // first, before trusting v. + int overflow = 0; + long v = PyLong_AsLongAndOverflow(value, &overflow); + if (overflow == 0 && v == -1 && PyErr_Occurred()) + return -1; // non-overflow conversion error; exception already set + if (overflow != 0 || v < INT_MIN || v > INT_MAX) + { + PyErr_SetString(PyExc_OverflowError, + "Python int is out of range for a c_int (32-bit) kernel argument"); + return -1; + } + *((int*)ptr) = (int)v; return sizeof(int); }; return; @@ -89,7 +108,19 @@ static void populate_feeders(PyTypeObject* target_t, PyTypeObject* source_t) { m_feeders[{target_t,source_t}] = [](void* ptr, PyObject* value) -> int { - *((int8_t*)ptr) = (int8_t)PyLong_AsLong(value); + // Same rationale as the c_int feeder above; here the slot is an + // 8-bit c_byte, so bound to INT8_MIN/INT8_MAX. + int overflow = 0; + long v = PyLong_AsLongAndOverflow(value, &overflow); + if (overflow == 0 && v == -1 && PyErr_Occurred()) + return -1; // non-overflow conversion error; exception already set + if (overflow != 0 || v < INT8_MIN || v > INT8_MAX) + { + PyErr_SetString(PyExc_OverflowError, + "Python int is out of range for a c_byte (8-bit) kernel argument"); + return -1; + } + *((int8_t*)ptr) = (int8_t)v; return sizeof(int8_t); }; return; diff --git a/cuda_bindings/cuda/bindings/_lib/param_packer.pxd b/cuda_bindings/cuda/bindings/_lib/param_packer.pxd index 1c0ad690be4..d1f84059db1 100644 --- a/cuda_bindings/cuda/bindings/_lib/param_packer.pxd +++ b/cuda_bindings/cuda/bindings/_lib/param_packer.pxd @@ -4,4 +4,4 @@ # Include "param_packer.h" so its contents get compiled into every # Cython extension module that depends on param_packer.pxd. cdef extern from "param_packer.h": - int feed(void* ptr, object o, object ct) + int feed(void* ptr, object o, object ct) except? -1 diff --git a/cuda_bindings/tests/test_kernelParams.py b/cuda_bindings/tests/test_kernelParams.py index 555d6a7284c..1e48d2d2fd6 100644 --- a/cuda_bindings/tests/test_kernelParams.py +++ b/cuda_bindings/tests/test_kernelParams.py @@ -800,3 +800,34 @@ def __init__(self, address, typestr): ASSERT_DRV(err) (err,) = cuda.cuModuleUnload(module) ASSERT_DRV(err) + + +def test_kernelParams_c_int_out_of_range_raises(device): + # #363: an out-of-range Python int for a c_int / c_byte kernel argument must + # raise instead of being silently truncated to fit the declared width. + kernelString = """\ + extern "C" __global__ void take_int(int i) {} + """ + module = common_nvrtc(kernelString, device) + err, kernel = cuda.cuModuleGetFunction(module, b"take_int") + ASSERT_DRV(err) + err, stream = cuda.cuStreamCreate(0) + ASSERT_DRV(err) + + # An in-range value still packs and launches fine. + (err,) = cuda.cuLaunchKernel(kernel, 1, 1, 1, 1, 1, 1, 0, stream, ((5,), (ctypes.c_int,)), 0) + ASSERT_DRV(err) + + # Out-of-range values now raise OverflowError during packing (previously the + # high bits were silently dropped, so the kernel saw a different value). + with pytest.raises(OverflowError): + cuda.cuLaunchKernel(kernel, 1, 1, 1, 1, 1, 1, 0, stream, ((2**32 + 5,), (ctypes.c_int,)), 0) + with pytest.raises(OverflowError): + cuda.cuLaunchKernel(kernel, 1, 1, 1, 1, 1, 1, 0, stream, ((200,), (ctypes.c_byte,)), 0) + + (err,) = cuda.cuStreamSynchronize(stream) + ASSERT_DRV(err) + (err,) = cuda.cuStreamDestroy(stream) + ASSERT_DRV(err) + (err,) = cuda.cuModuleUnload(module) + ASSERT_DRV(err)