diff --git a/cuda_bindings/cuda/bindings/_lib/param_packer.h b/cuda_bindings/cuda/bindings/_lib/param_packer.h index 160ef5f7c92..8d4833bb200 100644 --- a/cuda_bindings/cuda/bindings/_lib/param_packer.h +++ b/cuda_bindings/cuda/bindings/_lib/param_packer.h @@ -7,6 +7,28 @@ #include #include #include +#include +#include + +// PyLong_AsInt entered the public/stable CPython API in 3.13. cuda.bindings +// supports Python 3.10+, so provide a file-local backport for older builds. +// This is a copy of the CPython implementation; it is `static` (unlike the +// original) because this header is compiled into every extension module that +// includes it, mirroring the other helpers below. +#if PY_VERSION_HEX < 0x030D0000 +static int +PyLong_AsInt(PyObject *obj) +{ + int overflow; + long result = PyLong_AsLongAndOverflow(obj, &overflow); + if (overflow || result > INT_MAX || result < INT_MIN) { + PyErr_SetString(PyExc_OverflowError, + "Python int too large to convert to C int"); + return -1; + } + return (int)result; +} +#endif static PyObject* ctypes_module = nullptr; @@ -69,7 +91,13 @@ 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); + // PyLong_AsInt range-checks against the 32-bit int slot and raises + // OverflowError itself, so an out-of-range value is rejected rather + // than silently truncated. + int v = PyLong_AsInt(value); + if (v == -1 && PyErr_Occurred()) + return -1; + *((int*)ptr) = v; return sizeof(int); }; return; @@ -89,7 +117,23 @@ 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); + // c_byte is an 8-bit slot with no dedicated CPython converter, so + // range-check explicitly against INT8_MIN/INT8_MAX. AsLongAndOverflow's + // `overflow` only flags values outside `long` (64-bit on LP64), so a + // value in that range would be silently truncated by (int8_t)v without + // the explicit bounds check. When overflow!=0, v is the -1 sentinel + // (not the real value), so that case must be caught 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 < 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)