refactor: fully remove USE_NV_BINDING (#525)

Purge `USE_NV_BINDING` from the codebase.

Author: cpcloud

numba_cuda/numba/cuda/api.py

@@ -271,12 +271,8 @@ def open_ipc_array(handle, shape, dtype, strides=None, offset=0):
     # compute size
     size = np.prod(shape) * dtype.itemsize
     # manually recreate the IPC mem handle
-    if driver.USE_NV_BINDING:
-        driver_handle = driver.binding.CUipcMemHandle()
-        driver_handle.reserved = handle
-    else:
-        driver_handle = driver.drvapi.cu_ipc_mem_handle()
-        driver_handle.reserved[:] = handle
+    driver_handle = driver.binding.CUipcMemHandle()
+    driver_handle.reserved = handle
     # use *IpcHandle* to open the IPC memory
     ipchandle = driver.IpcHandle(None, driver_handle, size, offset=offset)
     yield ipchandle.open_array(

numba_cuda/numba/cuda/cudadrv/devicearray.py

@@ -112,10 +112,7 @@ def __init__(self, shape, strides, dtype, stream=0, gpu_data=None):
                 gpu_data = devices.get_context().memalloc(self.alloc_size)
         else:
             # Make NULL pointer for empty allocation
-            if _driver.USE_NV_BINDING:
-                null = _driver.binding.CUdeviceptr(0)
-            else:
-                null = c_void_p(0)
+            null = _driver.binding.CUdeviceptr(0)
             gpu_data = _driver.MemoryPointer(
                 context=devices.get_context(), pointer=null, size=0
             )

numba_cuda/numba/cuda/cudadrv/devices.py

@@ -18,7 +18,7 @@
 import threading
 from contextlib import contextmanager
 
-from .driver import driver, USE_NV_BINDING
+from .driver import driver
 
 
 class _DeviceList(object):
@@ -43,7 +43,7 @@ def __getitem__(self, devnum):
         """
         Returns the context manager for device *devnum*.
         """
-        if not isinstance(devnum, (int, slice)) and USE_NV_BINDING:
+        if not isinstance(devnum, (int, slice)):
             devnum = int(devnum)
         return self.lst[devnum]
 
@@ -146,8 +146,7 @@ def get_or_create_context(self, devnum):
             else:
                 return attached_ctx
         else:
-            if USE_NV_BINDING:
-                devnum = int(devnum)
+            devnum = int(devnum)
             return self._activate_context_for(devnum)
 
     def _get_or_create_context_uncached(self, devnum):

numba_cuda/numba/cuda/cudadrv/driver.py

@@ -63,10 +63,6 @@
     ObjectCode,
 )
 
-# For backwards compatibility: indicate that the NVIDIA CUDA Python bindings are
-# in use. Older code checks this flag to branch on binding-specific behavior.
-USE_NV_BINDING = True
-
 # There is no definition of the default stream in the Nvidia bindings (nor
 # is there at the C/C++ level), so we define it here so we don't need to
 # 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):
     try:
         fn(ptr, val, size, *varargs)
     except CudaAPIError as e:
-        invalid = (
-            binding.CUresult.CUDA_ERROR_INVALID_VALUE
-            if USE_NV_BINDING
-            else enums.CUDA_ERROR_INVALID_VALUE
-        )
+        invalid = binding.CUresult.CUDA_ERROR_INVALID_VALUE
         if (
             e.code == invalid
             and getattr(dst, "__cuda_memory__", False)

numba_cuda/numba/cuda/dispatcher.py

@@ -475,11 +475,7 @@ def launch(self, args, griddim, blockdim, stream=0, sharedmem=0):
         for t, v in zip(self.argument_types, args):
             self._prepare_args(t, v, stream, retr, kernelargs)
 
-        if driver.USE_NV_BINDING:
-            stream_handle = stream and stream.handle.value or 0
-        else:
-            zero_stream = None
-            stream_handle = stream and stream.handle or zero_stream
+        stream_handle = stream and stream.handle.value or 0
 
         # Invoke kernel
         driver.launch_kernel(
@@ -553,8 +549,7 @@ def _prepare_args(self, ty, val, stream, retr, kernelargs):
 
             ptr = driver.device_pointer(devary)
 
-            if driver.USE_NV_BINDING:
-                ptr = int(ptr)
+            ptr = int(ptr)
 
             data = ctypes.c_void_p(ptr)
 

numba_cuda/numba/cuda/memory_management/nrt.py

@@ -14,7 +14,6 @@
     _Linker,
     driver,
     launch_kernel,
-    USE_NV_BINDING,
     _have_nvjitlink,
 )
 from numba.cuda.cudadrv import devices
@@ -163,8 +162,7 @@ def allocate(self, stream=None):
         memsys_size = ctypes.c_uint64()
         ptr, nbytes = self._memsys_module.get_global_symbol("memsys_size")
         device_memsys_size = ptr.device_ctypes_pointer
-        if USE_NV_BINDING:
-            device_memsys_size = device_memsys_size.value
+        device_memsys_size = device_memsys_size.value
         driver.cuMemcpyDtoH(
             ctypes.addressof(memsys_size), device_memsys_size, nbytes
         )

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

@@ -62,8 +62,5 @@ def launch_kernel(*args, **kwargs):
     raise RuntimeError(msg)
 
 
-USE_NV_BINDING = False
-
-
 def _have_nvjitlink():
     return False

numba_cuda/numba/cuda/testing.py

@@ -299,10 +299,6 @@ def xfail_unless_cudasim(fn):
         return unittest.expectedFailure(fn)
 
 
-def skip_with_cuda_python(reason):
-    return unittest.skipIf(driver.USE_NV_BINDING, reason)
-
-
 def cudadevrt_missing():
     if config.ENABLE_CUDASIM:
         return False

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

@@ -2,7 +2,6 @@
 # SPDX-License-Identifier: BSD-2-Clause
 
 import numbers
-from ctypes import byref
 import weakref
 
 from numba import cuda
@@ -31,9 +30,6 @@ def test_gpus_cudevice_indexing(self):
         device_ids = [device.id for device in cuda.list_devices()]
         for device_id in device_ids:
             with cuda.gpus[device_id]:
-                # Check that the device is an integer if not using the CUDA
-                # Python bindings, otherwise it's a CUdevice object
-                assert isinstance(device_id, int) != driver.USE_NV_BINDING
                 self.assertEqual(cuda.gpus.current.id, device_id)
 
 
@@ -91,14 +87,9 @@ def tearDown(self):
     def test_attached_primary(self, extra_work=lambda: None):
         # Emulate primary context creation by 3rd party
         the_driver = driver.driver
-        if driver.USE_NV_BINDING:
-            dev = driver.binding.CUdevice(0)
-            binding_hctx = the_driver.cuDevicePrimaryCtxRetain(dev)
-            hctx = driver.drvapi.cu_context(int(binding_hctx))
-        else:
-            dev = 0
-            hctx = driver.drvapi.cu_context()
-            the_driver.cuDevicePrimaryCtxRetain(byref(hctx), dev)
+        dev = driver.binding.CUdevice(0)
+        binding_hctx = the_driver.cuDevicePrimaryCtxRetain(dev)
+        hctx = driver.drvapi.cu_context(int(binding_hctx))
         try:
             ctx = driver.Context(weakref.proxy(self), hctx)
             ctx.push()
@@ -115,33 +106,29 @@ def test_attached_primary(self, extra_work=lambda: None):
     def test_attached_non_primary(self):
         # Emulate non-primary context creation by 3rd party
         the_driver = driver.driver
-        if driver.USE_NV_BINDING:
-            flags = 0
-            dev = driver.binding.CUdevice(0)
-
-            result, version = driver.binding.cuDriverGetVersion()
-            self.assertEqual(
-                result,
-                driver.binding.CUresult.CUDA_SUCCESS,
-                "Error getting CUDA driver version",
-            )
-
-            # CUDA 13's cuCtxCreate has an optional parameter prepended. The
-            # version of cuCtxCreate in use depends on the cuda.bindings major
-            # version rather than the installed driver version on the machine
-            # we're running on.
-            from cuda import bindings
-
-            bindings_version = int(bindings.__version__.split(".")[0])
-            if bindings_version in (11, 12):
-                args = (flags, dev)
-            else:
-                args = (None, flags, dev)
-
-            hctx = the_driver.cuCtxCreate(*args)
+        flags = 0
+        dev = driver.binding.CUdevice(0)
+
+        result, version = driver.binding.cuDriverGetVersion()
+        self.assertEqual(
+            result,
+            driver.binding.CUresult.CUDA_SUCCESS,
+            "Error getting CUDA driver version",
+        )
+
+        # CUDA 13's cuCtxCreate has an optional parameter prepended. The
+        # version of cuCtxCreate in use depends on the cuda.bindings major
+        # version rather than the installed driver version on the machine
+        # we're running on.
+        from cuda import bindings
+
+        bindings_version = int(bindings.__version__.split(".")[0])
+        if bindings_version in (11, 12):
+            args = (flags, dev)
         else:
-            hctx = driver.drvapi.cu_context()
-            the_driver.cuCtxCreate(byref(hctx), 0, 0)
+            args = (None, flags, dev)
+
+        hctx = the_driver.cuCtxCreate(*args)
         try:
             cuda.current_context()
         except RuntimeError as e:

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

@@ -1,15 +1,15 @@
 # SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
 # SPDX-License-Identifier: BSD-2-Clause
 
-from ctypes import byref, c_int, sizeof
+from ctypes import c_int, sizeof
 
 from numba.cuda.cudadrv.driver import (
     host_to_device,
     device_to_host,
     driver,
     launch_kernel,
 )
-from numba.cuda.cudadrv import devices, drvapi, driver as _driver
+from numba.cuda.cudadrv import devices, driver as _driver
 from numba.cuda.testing import unittest, CUDATestCase
 from numba.cuda.testing import skip_on_cudasim
 
@@ -96,8 +96,7 @@ def test_cuda_driver_basic(self):
         ptr = memory.device_ctypes_pointer
         stream = 0
 
-        if _driver.USE_NV_BINDING:
-            stream = _driver.binding.CUstream(stream)
+        stream = _driver.binding.CUstream(stream)
 
         launch_kernel(
             function.handle,  # Kernel
@@ -133,8 +132,7 @@ def test_cuda_driver_stream_operations(self):
             ptr = memory.device_ctypes_pointer
 
             stream_handle = stream.handle
-            if _driver.USE_NV_BINDING:
-                stream_handle = stream_handle.value
+            stream_handle = stream_handle.value
 
             launch_kernel(
                 function.handle,  # Kernel
@@ -195,13 +193,8 @@ def test_cuda_driver_external_stream(self):
         # Test properties of a stream created from an external stream object.
         # We use the driver API directly to create a stream, to emulate an
         # external library creating a stream
-        if _driver.USE_NV_BINDING:
-            handle = driver.cuStreamCreate(0)
-            ptr = int(handle)
-        else:
-            handle = drvapi.cu_stream()
-            driver.cuStreamCreate(byref(handle), 0)
-            ptr = handle.value
+        handle = driver.cuStreamCreate(0)
+        ptr = int(handle)
         s = self.context.create_external_stream(ptr)
 
         self.assertIn("External CUDA stream", repr(s))