Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
8 changes: 2 additions & 6 deletions numba_cuda/numba/cuda/api.py
Original file line number Diff line number Diff line change
Expand Up @@ -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(
Expand Down
5 changes: 1 addition & 4 deletions numba_cuda/numba/cuda/cudadrv/devicearray.py
Original file line number Diff line number Diff line change
Expand Up @@ -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
)
Expand Down
7 changes: 3 additions & 4 deletions numba_cuda/numba/cuda/cudadrv/devices.py
Original file line number Diff line number Diff line change
Expand Up @@ -18,7 +18,7 @@
import threading
from contextlib import contextmanager

from .driver import driver, USE_NV_BINDING
from .driver import driver


class _DeviceList(object):
Expand All @@ -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]

Expand Down Expand Up @@ -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):
Expand Down
10 changes: 1 addition & 9 deletions numba_cuda/numba/cuda/cudadrv/driver.py
Original file line number Diff line number Diff line change
Expand Up @@ -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.
Expand Down Expand Up @@ -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)
Expand Down
9 changes: 2 additions & 7 deletions numba_cuda/numba/cuda/dispatcher.py
Original file line number Diff line number Diff line change
Expand Up @@ -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(
Expand Down Expand Up @@ -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)

Expand Down
4 changes: 1 addition & 3 deletions numba_cuda/numba/cuda/memory_management/nrt.py
Original file line number Diff line number Diff line change
Expand Up @@ -14,7 +14,6 @@
_Linker,
driver,
launch_kernel,
USE_NV_BINDING,
_have_nvjitlink,
)
from numba.cuda.cudadrv import devices
Expand Down Expand Up @@ -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
)
Expand Down
3 changes: 0 additions & 3 deletions numba_cuda/numba/cuda/simulator/cudadrv/driver.py
Original file line number Diff line number Diff line change
Expand Up @@ -62,8 +62,5 @@ def launch_kernel(*args, **kwargs):
raise RuntimeError(msg)


USE_NV_BINDING = False


def _have_nvjitlink():
return False
4 changes: 0 additions & 4 deletions numba_cuda/numba/cuda/testing.py
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
63 changes: 25 additions & 38 deletions numba_cuda/numba/cuda/tests/cudadrv/test_context_stack.py
Original file line number Diff line number Diff line change
Expand Up @@ -2,7 +2,6 @@
# SPDX-License-Identifier: BSD-2-Clause

import numbers
from ctypes import byref
import weakref

from numba import cuda
Expand Down Expand Up @@ -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)


Expand Down Expand Up @@ -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()
Expand All @@ -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:
Expand Down
19 changes: 6 additions & 13 deletions numba_cuda/numba/cuda/tests/cudadrv/test_cuda_driver.py
Original file line number Diff line number Diff line change
@@ -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

Expand Down Expand Up @@ -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
Expand Down Expand Up @@ -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
Expand Down Expand Up @@ -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))
Expand Down
5 changes: 1 addition & 4 deletions numba_cuda/numba/cuda/tests/cudadrv/test_cuda_memory.py
Original file line number Diff line number Diff line change
Expand Up @@ -54,10 +54,7 @@ def test_derived_pointer(self):
# Use MemoryPointer.view to create derived pointer

def handle_val(mem):
if driver.USE_NV_BINDING:
return int(mem.handle)
else:
return mem.handle.value
return int(mem.handle)

def check(m, offset):
# create view
Expand Down
13 changes: 3 additions & 10 deletions numba_cuda/numba/cuda/tests/cudadrv/test_managed_alloc.py
Original file line number Diff line number Diff line change
Expand Up @@ -2,8 +2,7 @@
# SPDX-License-Identifier: BSD-2-Clause

import numpy as np
from ctypes import byref, c_size_t
from numba.cuda.cudadrv.driver import device_memset, driver, USE_NV_BINDING
from numba.cuda.cudadrv.driver import device_memset, driver
from numba import cuda
from numba.cuda.testing import unittest, CUDATestCase
from numba.cuda.testing import skip_on_cudasim, skip_on_arm
Expand All @@ -22,14 +21,8 @@ def get_total_gpu_memory(self):
# We use a driver function to directly get the total GPU memory because
# an EMM plugin may report something different (or not implement
# get_memory_info at all).
if USE_NV_BINDING:
free, total = driver.cuMemGetInfo()
return total
else:
free = c_size_t()
total = c_size_t()
driver.cuMemGetInfo(byref(free), byref(total))
return total.value
free, total = driver.cuMemGetInfo()
return total

def skip_if_cc_major_lt(self, min_required, reason):
"""
Expand Down
Loading