Raise OverflowError on out-of-range c_int/c_byte kernel arguments

cuda_bindings/cuda/bindings/_lib/param_packer.h

@@ -7,6 +7,8 @@
 #include <functional>
 #include <stdexcept>
 #include <string>
+#include <climits>
+#include <cstdint>
 
 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;

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

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)