Raise OverflowError on out-of-range c_int/c_byte kernel arguments (#2402)

* 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) <noreply@anthropic.com>

* 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) <noreply@anthropic.com>

* param_packer: use PyLong_AsInt for c_int range check (backport for <3.13)

Address review feedback on #2402: replace the hand-rolled c_int bounds
check with CPython's PyLong_AsInt, which does exactly the INT_MIN/INT_MAX
range check we want and raises OverflowError itself. PyLong_AsInt only
became public in Python 3.13, so add a file-local, version-gated backport
(a copy of the CPython implementation, made `static` because this header
is compiled into every extension module that includes it).

The c_byte feeder has no CPython equivalent, so it keeps the explicit
PyLong_AsLongAndOverflow bounds check; its comment is now self-contained
rather than referencing the (now-changed) c_int feeder.

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

---------

Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

Author: rparolin

cuda_bindings/cuda/bindings/_lib/param_packer.h

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

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

@@ -782,3 +782,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)