Skip to content

Commit d8bd3a7

Browse files
cpcloudcpcloud
committed
refactor: fully remove USE_NV_BINDING 
1 parent 35978a9 commit d8bd3a7

File tree

14 files changed

+50
-179
lines changed

14 files changed

+50
-179
lines changed

‎numba_cuda/numba/cuda/api.py‎

Lines changed: 2 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -271,12 +271,8 @@ def open_ipc_array(handle, shape, dtype, strides=None, offset=0):
271271
# compute size
272272
size = np.prod(shape) * dtype.itemsize
273273
# manually recreate the IPC mem handle
274-
if driver.USE_NV_BINDING:
275-
driver_handle = driver.binding.CUipcMemHandle()
276-
driver_handle.reserved = handle
277-
else:
278-
driver_handle = driver.drvapi.cu_ipc_mem_handle()
279-
driver_handle.reserved[:] = handle
274+
driver_handle = driver.binding.CUipcMemHandle()
275+
driver_handle.reserved = handle
280276
# use *IpcHandle* to open the IPC memory
281277
ipchandle = driver.IpcHandle(None, driver_handle, size, offset=offset)
282278
yield ipchandle.open_array(

‎numba_cuda/numba/cuda/cudadrv/devicearray.py‎

Lines changed: 1 addition & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -112,10 +112,7 @@ def __init__(self, shape, strides, dtype, stream=0, gpu_data=None):
112112
gpu_data = devices.get_context().memalloc(self.alloc_size)
113113
else:
114114
# Make NULL pointer for empty allocation
115-
if _driver.USE_NV_BINDING:
116-
null = _driver.binding.CUdeviceptr(0)
117-
else:
118-
null = c_void_p(0)
115+
null = _driver.binding.CUdeviceptr(0)
119116
gpu_data = _driver.MemoryPointer(
120117
context=devices.get_context(), pointer=null, size=0
121118
)

‎numba_cuda/numba/cuda/cudadrv/devices.py‎

Lines changed: 3 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -18,7 +18,7 @@
1818
import threading
1919
from contextlib import contextmanager
2020

21-
from .driver import driver, USE_NV_BINDING
21+
from .driver import driver
2222

2323

2424
class _DeviceList(object):
@@ -43,7 +43,7 @@ def __getitem__(self, devnum):
4343
"""
4444
Returns the context manager for device *devnum*.
4545
"""
46-
if not isinstance(devnum, (int, slice)) and USE_NV_BINDING:
46+
if not isinstance(devnum, (int, slice)):
4747
devnum = int(devnum)
4848
return self.lst[devnum]
4949

@@ -146,8 +146,7 @@ def get_or_create_context(self, devnum):
146146
else:
147147
return attached_ctx
148148
else:
149-
if USE_NV_BINDING:
150-
devnum = int(devnum)
149+
devnum = int(devnum)
151150
return self._activate_context_for(devnum)
152151

153152
def _get_or_create_context_uncached(self, devnum):

‎numba_cuda/numba/cuda/cudadrv/driver.py‎

Lines changed: 1 addition & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -63,10 +63,6 @@
6363
ObjectCode,
6464
)
6565

66-
# For backwards compatibility: indicate that the NVIDIA CUDA Python bindings are
67-
# in use. Older code checks this flag to branch on binding-specific behavior.
68-
USE_NV_BINDING = True
69-
7066
# There is no definition of the default stream in the Nvidia bindings (nor
7167
# is there at the C/C++ level), so we define it here so we don't need to
7268
# use a magic number 0 in places where we want the default stream.
@@ -3176,11 +3172,7 @@ def device_memset(dst, val, size, stream=0):
31763172
try:
31773173
fn(ptr, val, size, *varargs)
31783174
except CudaAPIError as e:
3179-
invalid = (
3180-
binding.CUresult.CUDA_ERROR_INVALID_VALUE
3181-
if USE_NV_BINDING
3182-
else enums.CUDA_ERROR_INVALID_VALUE
3183-
)
3175+
invalid = binding.CUresult.CUDA_ERROR_INVALID_VALUE
31843176
if (
31853177
e.code == invalid
31863178
and getattr(dst, "__cuda_memory__", False)

‎numba_cuda/numba/cuda/dispatcher.py‎

Lines changed: 2 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -475,11 +475,7 @@ def launch(self, args, griddim, blockdim, stream=0, sharedmem=0):
475475
for t, v in zip(self.argument_types, args):
476476
self._prepare_args(t, v, stream, retr, kernelargs)
477477

478-
if driver.USE_NV_BINDING:
479-
stream_handle = stream and stream.handle.value or 0
480-
else:
481-
zero_stream = None
482-
stream_handle = stream and stream.handle or zero_stream
478+
stream_handle = stream and stream.handle.value or 0
483479

484480
# Invoke kernel
485481
driver.launch_kernel(
@@ -553,8 +549,7 @@ def _prepare_args(self, ty, val, stream, retr, kernelargs):
553549

554550
ptr = driver.device_pointer(devary)
555551

556-
if driver.USE_NV_BINDING:
557-
ptr = int(ptr)
552+
ptr = int(ptr)
558553

559554
data = ctypes.c_void_p(ptr)
560555

‎numba_cuda/numba/cuda/memory_management/nrt.py‎

Lines changed: 1 addition & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -14,7 +14,6 @@
1414
_Linker,
1515
driver,
1616
launch_kernel,
17-
USE_NV_BINDING,
1817
_have_nvjitlink,
1918
)
2019
from numba.cuda.cudadrv import devices
@@ -163,8 +162,7 @@ def allocate(self, stream=None):
163162
memsys_size = ctypes.c_uint64()
164163
ptr, nbytes = self._memsys_module.get_global_symbol("memsys_size")
165164
device_memsys_size = ptr.device_ctypes_pointer
166-
if USE_NV_BINDING:
167-
device_memsys_size = device_memsys_size.value
165+
device_memsys_size = device_memsys_size.value
168166
driver.cuMemcpyDtoH(
169167
ctypes.addressof(memsys_size), device_memsys_size, nbytes
170168
)

‎numba_cuda/numba/cuda/simulator/cudadrv/driver.py‎

Lines changed: 0 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -62,8 +62,5 @@ def launch_kernel(*args, **kwargs):
6262
raise RuntimeError(msg)
6363

6464

65-
USE_NV_BINDING = False
66-
67-
6865
def _have_nvjitlink():
6966
return False

‎numba_cuda/numba/cuda/testing.py‎

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -299,10 +299,6 @@ def xfail_unless_cudasim(fn):
299299
return unittest.expectedFailure(fn)
300300

301301

302-
def skip_with_cuda_python(reason):
303-
return unittest.skipIf(driver.USE_NV_BINDING, reason)
304-
305-
306302
def cudadevrt_missing():
307303
if config.ENABLE_CUDASIM:
308304
return False

‎numba_cuda/numba/cuda/tests/cudadrv/test_context_stack.py‎

Lines changed: 26 additions & 38 deletions
Original file line numberDiff line numberDiff line change
@@ -2,7 +2,6 @@
22
# SPDX-License-Identifier: BSD-2-Clause
33

44
import numbers
5-
from ctypes import byref
65
import weakref
76

87
from numba import cuda
@@ -31,9 +30,7 @@ def test_gpus_cudevice_indexing(self):
3130
device_ids = [device.id for device in cuda.list_devices()]
3231
for device_id in device_ids:
3332
with cuda.gpus[device_id]:
34-
# Check that the device is an integer if not using the CUDA
35-
# Python bindings, otherwise it's a CUdevice object
36-
assert isinstance(device_id, int) != driver.USE_NV_BINDING
33+
assert not isinstance(device_id, int)
3734
self.assertEqual(cuda.gpus.current.id, device_id)
3835

3936

@@ -91,14 +88,9 @@ def tearDown(self):
9188
def test_attached_primary(self, extra_work=lambda: None):
9289
# Emulate primary context creation by 3rd party
9390
the_driver = driver.driver
94-
if driver.USE_NV_BINDING:
95-
dev = driver.binding.CUdevice(0)
96-
binding_hctx = the_driver.cuDevicePrimaryCtxRetain(dev)
97-
hctx = driver.drvapi.cu_context(int(binding_hctx))
98-
else:
99-
dev = 0
100-
hctx = driver.drvapi.cu_context()
101-
the_driver.cuDevicePrimaryCtxRetain(byref(hctx), dev)
91+
dev = driver.binding.CUdevice(0)
92+
binding_hctx = the_driver.cuDevicePrimaryCtxRetain(dev)
93+
hctx = driver.drvapi.cu_context(int(binding_hctx))
10294
try:
10395
ctx = driver.Context(weakref.proxy(self), hctx)
10496
ctx.push()
@@ -115,33 +107,29 @@ def test_attached_primary(self, extra_work=lambda: None):
115107
def test_attached_non_primary(self):
116108
# Emulate non-primary context creation by 3rd party
117109
the_driver = driver.driver
118-
if driver.USE_NV_BINDING:
119-
flags = 0
120-
dev = driver.binding.CUdevice(0)
121-
122-
result, version = driver.binding.cuDriverGetVersion()
123-
self.assertEqual(
124-
result,
125-
driver.binding.CUresult.CUDA_SUCCESS,
126-
"Error getting CUDA driver version",
127-
)
128-
129-
# CUDA 13's cuCtxCreate has an optional parameter prepended. The
130-
# version of cuCtxCreate in use depends on the cuda.bindings major
131-
# version rather than the installed driver version on the machine
132-
# we're running on.
133-
from cuda import bindings
134-
135-
bindings_version = int(bindings.__version__.split(".")[0])
136-
if bindings_version in (11, 12):
137-
args = (flags, dev)
138-
else:
139-
args = (None, flags, dev)
140-
141-
hctx = the_driver.cuCtxCreate(*args)
110+
flags = 0
111+
dev = driver.binding.CUdevice(0)
112+
113+
result, version = driver.binding.cuDriverGetVersion()
114+
self.assertEqual(
115+
result,
116+
driver.binding.CUresult.CUDA_SUCCESS,
117+
"Error getting CUDA driver version",
118+
)
119+
120+
# CUDA 13's cuCtxCreate has an optional parameter prepended. The
121+
# version of cuCtxCreate in use depends on the cuda.bindings major
122+
# version rather than the installed driver version on the machine
123+
# we're running on.
124+
from cuda import bindings
125+
126+
bindings_version = int(bindings.__version__.split(".")[0])
127+
if bindings_version in (11, 12):
128+
args = (flags, dev)
142129
else:
143-
hctx = driver.drvapi.cu_context()
144-
the_driver.cuCtxCreate(byref(hctx), 0, 0)
130+
args = (None, flags, dev)
131+
132+
hctx = the_driver.cuCtxCreate(*args)
145133
try:
146134
cuda.current_context()
147135
except RuntimeError as e:

‎numba_cuda/numba/cuda/tests/cudadrv/test_cuda_driver.py‎

Lines changed: 6 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -1,15 +1,15 @@
11
# SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
22
# SPDX-License-Identifier: BSD-2-Clause
33

4-
from ctypes import byref, c_int, sizeof
4+
from ctypes import c_int, sizeof
55

66
from numba.cuda.cudadrv.driver import (
77
host_to_device,
88
device_to_host,
99
driver,
1010
launch_kernel,
1111
)
12-
from numba.cuda.cudadrv import devices, drvapi, driver as _driver
12+
from numba.cuda.cudadrv import devices, driver as _driver
1313
from numba.cuda.testing import unittest, CUDATestCase
1414
from numba.cuda.testing import skip_on_cudasim
1515

@@ -96,8 +96,7 @@ def test_cuda_driver_basic(self):
9696
ptr = memory.device_ctypes_pointer
9797
stream = 0
9898

99-
if _driver.USE_NV_BINDING:
100-
stream = _driver.binding.CUstream(stream)
99+
stream = _driver.binding.CUstream(stream)
101100

102101
launch_kernel(
103102
function.handle, # Kernel
@@ -133,8 +132,7 @@ def test_cuda_driver_stream_operations(self):
133132
ptr = memory.device_ctypes_pointer
134133

135134
stream_handle = stream.handle
136-
if _driver.USE_NV_BINDING:
137-
stream_handle = stream_handle.value
135+
stream_handle = stream_handle.value
138136

139137
launch_kernel(
140138
function.handle, # Kernel
@@ -195,13 +193,8 @@ def test_cuda_driver_external_stream(self):
195193
# Test properties of a stream created from an external stream object.
196194
# We use the driver API directly to create a stream, to emulate an
197195
# external library creating a stream
198-
if _driver.USE_NV_BINDING:
199-
handle = driver.cuStreamCreate(0)
200-
ptr = int(handle)
201-
else:
202-
handle = drvapi.cu_stream()
203-
driver.cuStreamCreate(byref(handle), 0)
204-
ptr = handle.value
196+
handle = driver.cuStreamCreate(0)
197+
ptr = int(handle)
205198
s = self.context.create_external_stream(ptr)
206199

207200
self.assertIn("External CUDA stream", repr(s))

0 commit comments

Comments
 (0)