From e2ef9963e5b7039fbb835b9981e161b92ce26246 Mon Sep 17 00:00:00 2001 From: Rob Parolin Date: Tue, 21 Jul 2026 15:58:44 -0700 Subject: [PATCH 1/2] Raise OverflowError on out-of-range c_int/c_byte kernel arguments An out-of-range Python int passed as a c_int/c_byte kernel argument was silently narrowed (e.g. 2**32+5 -> 5). The feeders now range-check via PyLong_AsLongAndOverflow + an explicit 32-bit/8-bit bounds check and raise OverflowError; feed() is declared 'except? -1' so Cython propagates the exception instead of swallowing it. Adds a GPU regression test. Addresses Glasswing finding V9.1; behavior reviewed and shaped by @leofang (single code path via PyLong_AsLongAndOverflow). Co-Authored-By: Claude Opus 4.8 (1M context) --- .../cuda/bindings/_lib/param_packer.h | 29 +++++++++++++++-- .../cuda/bindings/_lib/param_packer.pxd | 2 +- cuda_bindings/tests/test_kernelParams.py | 31 +++++++++++++++++++ 3 files changed, 59 insertions(+), 3 deletions(-) diff --git a/cuda_bindings/cuda/bindings/_lib/param_packer.h b/cuda_bindings/cuda/bindings/_lib/param_packer.h index 160ef5f7c92..7161094a447 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,20 @@ 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); + // One code path across all supported Pythons: AsLongAndOverflow + // flags values outside C long via `overflow` (without setting an + // exception), then we bounds-check the 32-bit int range. + 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 +104,17 @@ 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); + 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) From d94f7b06ee3d2a9c5adfda8afa92dfce0c43a382 Mon Sep 17 00:00:00 2001 From: Rob Parolin Date: Tue, 21 Jul 2026 16:16:08 -0700 Subject: [PATCH 2/2] param_packer: clarify why c_int/c_byte bound to int range, not long A reviewer asked why the manual check uses INT_MIN/INT_MAX rather than the function's long-based overflow flag. Add a comment: the target slot is 32-bit (8-bit for c_byte), and long is 64-bit on LP64, so overflow alone misses 2**31..2**63 -- the explicit int-range check catches those. Comment-only. Co-Authored-By: Claude Opus 4.8 (1M context) --- cuda_bindings/cuda/bindings/_lib/param_packer.h | 12 +++++++++--- 1 file changed, 9 insertions(+), 3 deletions(-) diff --git a/cuda_bindings/cuda/bindings/_lib/param_packer.h b/cuda_bindings/cuda/bindings/_lib/param_packer.h index 7161094a447..5cbf6544fb4 100644 --- a/cuda_bindings/cuda/bindings/_lib/param_packer.h +++ b/cuda_bindings/cuda/bindings/_lib/param_packer.h @@ -71,9 +71,13 @@ static void populate_feeders(PyTypeObject* target_t, PyTypeObject* source_t) { m_feeders[{target_t,source_t}] = [](void* ptr, PyObject* value) -> int { - // One code path across all supported Pythons: AsLongAndOverflow - // flags values outside C long via `overflow` (without setting an - // exception), then we bounds-check the 32-bit int range. + // 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()) @@ -104,6 +108,8 @@ static void populate_feeders(PyTypeObject* target_t, PyTypeObject* source_t) { m_feeders[{target_t,source_t}] = [](void* ptr, PyObject* value) -> int { + // 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())