Merge branch 'main' into fix_2183

Author: lijinf2

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)

cuda_core/cuda/core/_device_resources.pxd

@@ -2,8 +2,6 @@
 #
 # SPDX-License-Identifier: Apache-2.0
 
-cimport cython
-
 from cuda.bindings cimport cydriver
 from cuda.core._resource_handles cimport ContextHandle, GreenCtxHandle
 
@@ -17,7 +15,6 @@ cdef class SMResource:
         unsigned int _flags
         bint _is_usable
         object __weakref__
-        cython.pymutex _split_mutex
 
     @staticmethod
     cdef SMResource _from_dev_resource(cydriver.CUdevResource res, int device_id)

cuda_core/cuda/core/_device_resources.pyx

@@ -520,11 +520,10 @@ cdef class SMResource:
         )
         _resolve_group_count(opts)
         _check_green_ctx_support()
-        with self._split_mutex:
-            if _can_use_structured_sm_split():
-                return _split_with_general_api(self, opts, dry_run)
-            # SplitByCount requires the same 12.4+ as green ctx support (already checked above)
-            return _split_with_count_api(self, opts, dry_run)
+        if _can_use_structured_sm_split():
+            return _split_with_general_api(self, opts, dry_run)
+        # SplitByCount requires the same 12.4+ as green ctx support (already checked above)
+        return _split_with_count_api(self, opts, dry_run)
 
 
 cdef class WorkqueueResource:

cuda_core/cuda/core/_module.pyi

@@ -458,6 +458,20 @@ class ObjectCode:
 
         """
 
+    def get_module(self) -> object:
+        """Return a context-dependent :obj:`~driver.CUmodule` for legacy interop.
+
+        Bridges the native :obj:`~driver.CUlibrary` (see :attr:`handle`) to a
+        ``CUmodule`` via ``cuLibraryGetModule``, for use with legacy driver APIs
+        that only accept ``CUmodule``.
+
+        Returns
+        -------
+        :obj:`~driver.CUmodule`
+            Module handle for the current CUDA context, suitable for legacy
+            driver APIs that accept ``CUmodule``.
+        """
+
     @property
     def code(self) -> CodeTypeT:
         """Return the underlying code object."""
@@ -476,7 +490,10 @@ class ObjectCode:
 
     @property
     def handle(self) -> object:
-        """Return the underlying handle object.
+        """Return the native, context-independent :obj:`~driver.CUlibrary` handle.
+
+        Used by ``cuda.core`` and newer driver library APIs. For legacy APIs
+        that only accept a ``CUmodule``, use :meth:`get_module` instead.
 
         .. caution::
 

cuda_core/cuda/core/_module.pyx

@@ -6,6 +6,7 @@ from __future__ import annotations
 
 cimport cython
 from libc.stddef cimport size_t
+from libc.stdint cimport intptr_t
 
 from collections import namedtuple
 from os import fsencode, fspath, PathLike
@@ -796,6 +797,25 @@ cdef class ObjectCode:
             HANDLE_RETURN(get_last_error())
         return Kernel._from_handle(h_kernel)
 
+    def get_module(self) -> object:
+        """Return a context-dependent :obj:`~driver.CUmodule` for legacy interop.
+
+        Bridges the native :obj:`~driver.CUlibrary` (see :attr:`handle`) to a
+        ``CUmodule`` via ``cuLibraryGetModule``, for use with legacy driver APIs
+        that only accept ``CUmodule``.
+
+        Returns
+        -------
+        :obj:`~driver.CUmodule`
+            Module handle for the current CUDA context, suitable for legacy
+            driver APIs that accept ``CUmodule``.
+        """
+        self._lazy_load_module()
+        cdef cydriver.CUmodule mod
+        with nogil:
+            HANDLE_RETURN(cydriver.cuLibraryGetModule(&mod, as_cu(self._h_library)))
+        return driver.CUmodule(<intptr_t>mod)
+
     @property
     def code(self) -> CodeTypeT:
         """Return the underlying code object."""
@@ -818,7 +838,10 @@ cdef class ObjectCode:
 
     @property
     def handle(self) -> object:
-        """Return the underlying handle object.
+        """Return the native, context-independent :obj:`~driver.CUlibrary` handle.
+
+        Used by ``cuda.core`` and newer driver library APIs. For legacy APIs
+        that only accept a ``CUmodule``, use :meth:`get_module` instead.
 
         .. caution::
 

cuda_core/tests/conftest.py

@@ -131,6 +131,18 @@ def wrapper(*args, **kwargs):
                 kwargs["mempool_device_x2"] = _mempool_device_impl(2)
             if "mempool_device_x3" in kwargs:
                 kwargs["mempool_device_x3"] = _mempool_device_impl(3)
+
+            # These are used by test_green_context.py.  The original fixtures include
+            # pytest.skip() but that should have correctly fired by this time.
+            if "sm_resource" in kwargs:
+                kwargs["sm_resource"] = device.resources.sm
+            if "wq_resource" in kwargs:
+                kwargs["wq_resource"] = device.resources.workqueue
+            if "green_ctx" in kwargs:
+                from cuda.core import ContextOptions, SMResourceOptions
+
+                groups, _ = device.resources.sm.split(SMResourceOptions(count=None))
+                kwargs["green_ctx"] = device.create_context(ContextOptions(resources=[groups[0]]))
             return func(*args, **kwargs)
 
     wrapper._cuda_core_worker_cuda_wrapped = True

cuda_core/tests/graph/test_graph_definition_lifetime.py

@@ -611,48 +611,97 @@ def test_user_object_cleanup_is_coalesced_on_python_thread(init_cuda):
 
 
 @pytest.mark.agent_authored(model="gpt-5.6")
-def test_pending_call_queue_saturation_preserves_cleanup(init_cuda):
+def test_pending_call_queue_saturation_preserves_cleanup(tmp_path):
     """A full CPython queue neither strands nor mis-threads cleanup."""
-    pending_callback_type = ctypes.CFUNCTYPE(ctypes.c_int, ctypes.c_void_p)
-    add_pending_call = ctypes.pythonapi.Py_AddPendingCall
-    add_pending_call.argtypes = [pending_callback_type, ctypes.c_void_p]
-    add_pending_call.restype = ctypes.c_int
+    code = f"timeout = {_FINALIZE_TIMEOUT!r}\n" + textwrap.dedent(
+        """
+        import ctypes
+        import gc
+        import threading
+        import time
+
+        from cuda.core import Device
+        from cuda.core.graph import GraphDefinition
 
-    @pending_callback_type
-    def noop_pending_call(_):
-        return 0
+        class ThreadRecordingCallback:
+            def __init__(self, finalized_threads):
+                self.finalized_threads = finalized_threads
 
-    finalized_threads = []
-    main_thread = threading.get_ident()
-    first_callback = _ThreadRecordingCallback(finalized_threads)
-    first_graph = GraphDefinition()
-    first_graph.callback(first_callback)
-    graph_holder = [first_graph]
-    worker_done = threading.Event()
-    queue_was_full = []
+            def __call__(self):
+                pass
+
+            def __del__(self):
+                self.finalized_threads.append(threading.get_ident())
+
+        def wait_until(predicate):
+            deadline = time.monotonic() + timeout
+            while True:
+                gc.collect()
+                if predicate():
+                    return
+                if time.monotonic() >= deadline:
+                    break
+                time.sleep(0)
+                time.sleep(0.02)
+            raise AssertionError(f"condition not satisfied within {timeout}s")
+
+        Device(0).set_current()
+
+        pending_callback_type = ctypes.CFUNCTYPE(ctypes.c_int, ctypes.c_void_p)
+        add_pending_call = ctypes.pythonapi.Py_AddPendingCall
+        add_pending_call.argtypes = [pending_callback_type, ctypes.c_void_p]
+        add_pending_call.restype = ctypes.c_int
+        make_pending_calls = ctypes.pythonapi.Py_MakePendingCalls
+        make_pending_calls.argtypes = []
+        make_pending_calls.restype = ctypes.c_int
+
+        @pending_callback_type
+        def noop_pending_call(_):
+            return 0
+
+        finalized_threads = []
+        main_thread = threading.get_ident()
+        first_callback = ThreadRecordingCallback(finalized_threads)
+        first_graph = GraphDefinition()
+        first_graph.callback(first_callback)
+        graph_holder = [first_graph]
+        worker_done = threading.Event()
+        queue_was_full = []
+
+        del first_callback, first_graph
+
+        def fill_queue_and_destroy():
+            while add_pending_call(noop_pending_call, None) == 0:
+                pass
+            queue_was_full.append(True)
+            graph_holder.clear()
+            worker_done.set()
 
-    del first_callback, first_graph
+        worker = threading.Thread(target=fill_queue_and_destroy)
+        worker.start()
+        assert worker_done.wait(timeout=5)
+        worker.join()
+        assert queue_was_full == [True]
 
-    def fill_queue_and_destroy():
-        while add_pending_call(noop_pending_call, None) == 0:
-            pass
-        queue_was_full.append(True)
-        graph_holder.clear()
-        worker_done.set()
-
-    worker = threading.Thread(target=fill_queue_and_destroy)
-    worker.start()
-    assert worker_done.wait(timeout=5)
-    worker.join()
-    assert queue_was_full == [True]
-
-    # A later safe cuda-core close retries after the main thread has had an
-    # opportunity to drain the foreign pending calls.
-    retry_builder = Device().create_graph_builder()
-    retry_builder.close()
-
-    _wait_until(lambda: len(finalized_threads) == 1)
-    assert set(finalized_threads) == {main_thread}
+        # Free space before the cuda-core entry that retries cleanup scheduling.
+        assert make_pending_calls() == 0
+
+        retry_builder = Device().create_graph_builder()
+        retry_builder.close()
+
+        wait_until(lambda: len(finalized_threads) == 1)
+        assert set(finalized_threads) == {main_thread}
+        """
+    )
+    result = subprocess.run(  # noqa: S603 - controlled interpreter probe
+        [sys.executable, "-c", code],
+        capture_output=True,
+        text=True,
+        timeout=60,
+        # Isolate the process-global pending-call queue from parallel tests.
+        cwd=tmp_path,
+    )
+    assert result.returncode == 0, result.stderr
 
 
 @pytest.mark.agent_authored(model="gpt-5.6")

cuda_core/tests/test_green_context.py

@@ -41,6 +41,8 @@
 # ---------------------------------------------------------------------------
 # Fixtures
 # ---------------------------------------------------------------------------
+# Note that the following fixtures (except fill_kernel) require per-thread setup
+# and are currently special cased to work with pytest-run-parallel in conftest.
 
 
 # Resource queries (dev.resources.sm, dev.resources.workqueue) can fail in