Skip to content

Commit 565479b

Browse files
mdboommdboom
committed
Merge remote-tracking branch 'upstream/main' into test/api-docs-consistency-2326
2 parents 93e6d48 + 0d22cb4 commit 565479b

10 files changed

Lines changed: 153 additions & 13 deletions

File tree

‎cuda_bindings/cuda/bindings/_lib/param_packer.h‎

Lines changed: 46 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -7,6 +7,28 @@
77
#include <functional>
88
#include <stdexcept>
99
#include <string>
10+
#include <climits>
11+
#include <cstdint>
12+
13+
// PyLong_AsInt entered the public/stable CPython API in 3.13. cuda.bindings
14+
// supports Python 3.10+, so provide a file-local backport for older builds.
15+
// This is a copy of the CPython implementation; it is `static` (unlike the
16+
// original) because this header is compiled into every extension module that
17+
// includes it, mirroring the other helpers below.
18+
#if PY_VERSION_HEX < 0x030D0000
19+
static int
20+
PyLong_AsInt(PyObject *obj)
21+
{
22+
int overflow;
23+
long result = PyLong_AsLongAndOverflow(obj, &overflow);
24+
if (overflow || result > INT_MAX || result < INT_MIN) {
25+
PyErr_SetString(PyExc_OverflowError,
26+
"Python int too large to convert to C int");
27+
return -1;
28+
}
29+
return (int)result;
30+
}
31+
#endif
1032

1133
static PyObject* ctypes_module = nullptr;
1234

@@ -69,7 +91,13 @@ static void populate_feeders(PyTypeObject* target_t, PyTypeObject* source_t)
6991
{
7092
m_feeders[{target_t,source_t}] = [](void* ptr, PyObject* value) -> int
7193
{
72-
*((int*)ptr) = (int)PyLong_AsLong(value);
94+
// PyLong_AsInt range-checks against the 32-bit int slot and raises
95+
// OverflowError itself, so an out-of-range value is rejected rather
96+
// than silently truncated.
97+
int v = PyLong_AsInt(value);
98+
if (v == -1 && PyErr_Occurred())
99+
return -1;
100+
*((int*)ptr) = v;
73101
return sizeof(int);
74102
};
75103
return;
@@ -89,7 +117,23 @@ static void populate_feeders(PyTypeObject* target_t, PyTypeObject* source_t)
89117
{
90118
m_feeders[{target_t,source_t}] = [](void* ptr, PyObject* value) -> int
91119
{
92-
*((int8_t*)ptr) = (int8_t)PyLong_AsLong(value);
120+
// c_byte is an 8-bit slot with no dedicated CPython converter, so
121+
// range-check explicitly against INT8_MIN/INT8_MAX. AsLongAndOverflow's
122+
// `overflow` only flags values outside `long` (64-bit on LP64), so a
123+
// value in that range would be silently truncated by (int8_t)v without
124+
// the explicit bounds check. When overflow!=0, v is the -1 sentinel
125+
// (not the real value), so that case must be caught before trusting v.
126+
int overflow = 0;
127+
long v = PyLong_AsLongAndOverflow(value, &overflow);
128+
if (overflow == 0 && v == -1 && PyErr_Occurred())
129+
return -1; // non-overflow conversion error; exception already set
130+
if (overflow != 0 || v < INT8_MIN || v > INT8_MAX)
131+
{
132+
PyErr_SetString(PyExc_OverflowError,
133+
"Python int is out of range for a c_byte (8-bit) kernel argument");
134+
return -1;
135+
}
136+
*((int8_t*)ptr) = (int8_t)v;
93137
return sizeof(int8_t);
94138
};
95139
return;

‎cuda_bindings/cuda/bindings/_lib/param_packer.pxd‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -4,4 +4,4 @@
44
# Include "param_packer.h" so its contents get compiled into every
55
# Cython extension module that depends on param_packer.pxd.
66
cdef extern from "param_packer.h":
7-
int feed(void* ptr, object o, object ct)
7+
int feed(void* ptr, object o, object ct) except? -1

‎cuda_bindings/tests/test_kernelParams.py‎

Lines changed: 31 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -782,3 +782,34 @@ def __init__(self, address, typestr):
782782
ASSERT_DRV(err)
783783
(err,) = cuda.cuModuleUnload(module)
784784
ASSERT_DRV(err)
785+
786+
787+
def test_kernelParams_c_int_out_of_range_raises(device):
788+
# #363: an out-of-range Python int for a c_int / c_byte kernel argument must
789+
# raise instead of being silently truncated to fit the declared width.
790+
kernelString = """\
791+
extern "C" __global__ void take_int(int i) {}
792+
"""
793+
module = common_nvrtc(kernelString, device)
794+
err, kernel = cuda.cuModuleGetFunction(module, b"take_int")
795+
ASSERT_DRV(err)
796+
err, stream = cuda.cuStreamCreate(0)
797+
ASSERT_DRV(err)
798+
799+
# An in-range value still packs and launches fine.
800+
(err,) = cuda.cuLaunchKernel(kernel, 1, 1, 1, 1, 1, 1, 0, stream, ((5,), (ctypes.c_int,)), 0)
801+
ASSERT_DRV(err)
802+
803+
# Out-of-range values now raise OverflowError during packing (previously the
804+
# high bits were silently dropped, so the kernel saw a different value).
805+
with pytest.raises(OverflowError):
806+
cuda.cuLaunchKernel(kernel, 1, 1, 1, 1, 1, 1, 0, stream, ((2**32 + 5,), (ctypes.c_int,)), 0)
807+
with pytest.raises(OverflowError):
808+
cuda.cuLaunchKernel(kernel, 1, 1, 1, 1, 1, 1, 0, stream, ((200,), (ctypes.c_byte,)), 0)
809+
810+
(err,) = cuda.cuStreamSynchronize(stream)
811+
ASSERT_DRV(err)
812+
(err,) = cuda.cuStreamDestroy(stream)
813+
ASSERT_DRV(err)
814+
(err,) = cuda.cuModuleUnload(module)
815+
ASSERT_DRV(err)

‎cuda_core/cuda/core/_device_resources.pxd‎

Lines changed: 0 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -2,8 +2,6 @@
22
#
33
# SPDX-License-Identifier: Apache-2.0
44

5-
cimport cython
6-
75
from cuda.bindings cimport cydriver
86
from cuda.core._resource_handles cimport ContextHandle, GreenCtxHandle
97

@@ -17,7 +15,6 @@ cdef class SMResource:
1715
unsigned int _flags
1816
bint _is_usable
1917
object __weakref__
20-
cython.pymutex _split_mutex
2118

2219
@staticmethod
2320
cdef SMResource _from_dev_resource(cydriver.CUdevResource res, int device_id)

‎cuda_core/cuda/core/_device_resources.pyx‎

Lines changed: 4 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -520,11 +520,10 @@ cdef class SMResource:
520520
)
521521
_resolve_group_count(opts)
522522
_check_green_ctx_support()
523-
with self._split_mutex:
524-
if _can_use_structured_sm_split():
525-
return _split_with_general_api(self, opts, dry_run)
526-
# SplitByCount requires the same 12.4+ as green ctx support (already checked above)
527-
return _split_with_count_api(self, opts, dry_run)
523+
if _can_use_structured_sm_split():
524+
return _split_with_general_api(self, opts, dry_run)
525+
# SplitByCount requires the same 12.4+ as green ctx support (already checked above)
526+
return _split_with_count_api(self, opts, dry_run)
528527

529528

530529
cdef class WorkqueueResource:

‎cuda_core/cuda/core/_module.pyi‎

Lines changed: 18 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -458,6 +458,20 @@ class ObjectCode:
458458
459459
"""
460460

461+
def get_module(self) -> object:
462+
"""Return a context-dependent :obj:`~driver.CUmodule` for legacy interop.
463+
464+
Bridges the native :obj:`~driver.CUlibrary` (see :attr:`handle`) to a
465+
``CUmodule`` via ``cuLibraryGetModule``, for use with legacy driver APIs
466+
that only accept ``CUmodule``.
467+
468+
Returns
469+
-------
470+
:obj:`~driver.CUmodule`
471+
Module handle for the current CUDA context, suitable for legacy
472+
driver APIs that accept ``CUmodule``.
473+
"""
474+
461475
@property
462476
def code(self) -> CodeTypeT:
463477
"""Return the underlying code object."""
@@ -476,7 +490,10 @@ class ObjectCode:
476490

477491
@property
478492
def handle(self) -> object:
479-
"""Return the underlying handle object.
493+
"""Return the native, context-independent :obj:`~driver.CUlibrary` handle.
494+
495+
Used by ``cuda.core`` and newer driver library APIs. For legacy APIs
496+
that only accept a ``CUmodule``, use :meth:`get_module` instead.
480497
481498
.. caution::
482499

‎cuda_core/cuda/core/_module.pyx‎

Lines changed: 24 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -6,6 +6,7 @@ from __future__ import annotations
66

77
cimport cython
88
from libc.stddef cimport size_t
9+
from libc.stdint cimport intptr_t
910

1011
from collections import namedtuple
1112
from os import fsencode, fspath, PathLike
@@ -796,6 +797,25 @@ cdef class ObjectCode:
796797
HANDLE_RETURN(get_last_error())
797798
return Kernel._from_handle(h_kernel)
798799

800+
def get_module(self) -> object:
801+
"""Return a context-dependent :obj:`~driver.CUmodule` for legacy interop.
802+
803+
Bridges the native :obj:`~driver.CUlibrary` (see :attr:`handle`) to a
804+
``CUmodule`` via ``cuLibraryGetModule``, for use with legacy driver APIs
805+
that only accept ``CUmodule``.
806+
807+
Returns
808+
-------
809+
:obj:`~driver.CUmodule`
810+
Module handle for the current CUDA context, suitable for legacy
811+
driver APIs that accept ``CUmodule``.
812+
"""
813+
self._lazy_load_module()
814+
cdef cydriver.CUmodule mod
815+
with nogil:
816+
HANDLE_RETURN(cydriver.cuLibraryGetModule(&mod, as_cu(self._h_library)))
817+
return driver.CUmodule(<intptr_t>mod)
818+
799819
@property
800820
def code(self) -> CodeTypeT:
801821
"""Return the underlying code object."""
@@ -818,7 +838,10 @@ cdef class ObjectCode:
818838

819839
@property
820840
def handle(self) -> object:
821-
"""Return the underlying handle object.
841+
"""Return the native, context-independent :obj:`~driver.CUlibrary` handle.
842+
843+
Used by ``cuda.core`` and newer driver library APIs. For legacy APIs
844+
that only accept a ``CUmodule``, use :meth:`get_module` instead.
822845

823846
.. caution::
824847

‎cuda_core/tests/conftest.py‎

Lines changed: 12 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -131,6 +131,18 @@ def wrapper(*args, **kwargs):
131131
kwargs["mempool_device_x2"] = _mempool_device_impl(2)
132132
if "mempool_device_x3" in kwargs:
133133
kwargs["mempool_device_x3"] = _mempool_device_impl(3)
134+
135+
# These are used by test_green_context.py. The original fixtures include
136+
# pytest.skip() but that should have correctly fired by this time.
137+
if "sm_resource" in kwargs:
138+
kwargs["sm_resource"] = device.resources.sm
139+
if "wq_resource" in kwargs:
140+
kwargs["wq_resource"] = device.resources.workqueue
141+
if "green_ctx" in kwargs:
142+
from cuda.core import ContextOptions, SMResourceOptions
143+
144+
groups, _ = device.resources.sm.split(SMResourceOptions(count=None))
145+
kwargs["green_ctx"] = device.create_context(ContextOptions(resources=[groups[0]]))
134146
return func(*args, **kwargs)
135147

136148
wrapper._cuda_core_worker_cuda_wrapped = True

‎cuda_core/tests/test_green_context.py‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -41,6 +41,8 @@
4141
# ---------------------------------------------------------------------------
4242
# Fixtures
4343
# ---------------------------------------------------------------------------
44+
# Note that the following fixtures (except fill_kernel) require per-thread setup
45+
# and are currently special cased to work with pytest-run-parallel in conftest.
4446

4547

4648
# Resource queries (dev.resources.sm, dev.resources.workqueue) can fail in

‎cuda_core/tests/test_module.py‎

Lines changed: 15 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -282,6 +282,21 @@ def test_get_kernel(init_cuda):
282282
assert object_code.get_kernel(b"ABC").handle is not None
283283

284284

285+
def test_object_code_get_module_for_legacy_integration(init_cuda):
286+
src = """
287+
extern "C" __global__ void ABC() { }
288+
extern "C" __global__ void DEF() { }
289+
"""
290+
object_code = Program(src, "c++").compile("cubin")
291+
292+
# Bridge: CUlibrary (new) → CUmodule (legacy)
293+
module = object_code.get_module()
294+
295+
# Legacy module-only API consumes it directly
296+
count = handle_return(driver.cuModuleGetFunctionCount(module))
297+
assert count == 2
298+
299+
285300
@pytest.mark.parametrize(
286301
"attr, expected_type",
287302
[

0 commit comments

Comments
 (0)