Merge branch 'main' into vk/typing

Author: VijayKandiah

.github/workflows/pr.yaml

@@ -19,13 +19,11 @@ jobs:
       - compute-matrix
       - build-conda
       - test-conda
-      - test-conda-ctypes-binding
       - test-simulator
       - build-wheels
       - build-wheels-windows
       - test-wheels-windows
       - test-wheels
-      - test-wheels-ctypes-binding
       - test-wheels-deps-wheels
       - test-thirdparty
       - build-docs
@@ -80,18 +78,6 @@ jobs:
       script: "ci/test_conda.sh"
       run_codecov: false
       matrix: ${{ needs.compute-matrix.outputs.TEST_MATRIX }}
-  test-conda-ctypes-binding:
-    needs:
-      - build-conda
-      - compute-matrix
-    uses: ./.github/workflows/conda-python-tests.yaml
-    with:
-      build_type: pull-request
-      script: "ci/test_conda_ctypes_binding.sh"
-      run_codecov: false
-      # This selects "ARCH=amd64 and CUDA >=12, with the latest supported Python for each CUDA major version".
-      matrix: ${{ needs.compute-matrix.outputs.TEST_MATRIX }}
-      matrix_filter: map(select(.ARCH == "amd64" and (.CUDA_VER | split(".") | .[0] | tonumber >= 12))) | group_by(.CUDA_VER|split(".")|map(tonumber)|.[0]) | map(max_by([(.PY_VER|split(".")|map(tonumber)), (.CUDA_VER|split(".")|map(tonumber))]))
   test-simulator:
     needs:
       - build-conda
@@ -124,17 +110,6 @@ jobs:
       build_type: pull-request
       script: "ci/test_wheel.sh false"
       matrix: ${{ needs.compute-matrix.outputs.TEST_MATRIX }}
-  test-wheels-ctypes-binding:
-    needs:
-      - build-wheels
-      - compute-matrix
-    uses: ./.github/workflows/wheels-test.yaml
-    with:
-      build_type: pull-request
-      script: "ci/test_wheel_ctypes_binding.sh"
-      # This selects "ARCH=amd64 and CUDA >=12, with the latest supported Python for each CUDA major version".
-      matrix: ${{ needs.compute-matrix.outputs.TEST_MATRIX }}
-      matrix_filter: map(select(.ARCH == "amd64" and (.CUDA_VER | split(".") | .[0] | tonumber >= 12))) | group_by(.CUDA_VER|split(".")|map(tonumber)|.[0]) | map(max_by([(.PY_VER|split(".")|map(tonumber)), (.CUDA_VER|split(".")|map(tonumber))]))
   test-wheels-deps-wheels:
     needs:
       - build-wheels

ci/test_conda_ctypes_binding.sh

@@ -103,12 +103,11 @@ target.
    Enable warnings if a kernel is launched with host memory which forces a copy to and
    from the device. This option is on by default (default value is 1).
 
-.. envvar:: NUMBA_CUDA_USE_NVIDIA_BINDING
+.. note::
 
-   When set to 1, Numba will attempt to use the `NVIDIA CUDA Python binding
-   <https://nvidia.github.io/cuda-python/>`_ to make calls to the driver API
-   instead of using its own ctypes binding. This defaults to 1 (on). Set to
-   0 to use the ctypes bindings.
+   Numba-CUDA always uses the NVIDIA CUDA Python bindings. The legacy ctypes
+   bindings and the ``NUMBA_CUDA_USE_NVIDIA_BINDING`` environment variable have
+   been removed.
 
 .. envvar:: NUMBA_CUDA_INCLUDE_PATH
 

ci/test_wheel_ctypes_binding.sh

@@ -5,25 +5,22 @@
 CUDA Bindings
 =============
 
-Numba supports two bindings to the CUDA Driver APIs: its own internal bindings
-based on ctypes, and the official `NVIDIA CUDA Python bindings
-<https://nvidia.github.io/cuda-python/>`_. Functionality is equivalent between
-the two bindings.
-
-The internal bindings are used by default. If the NVIDIA bindings are installed,
-then they can be used by setting the environment variable
-``NUMBA_CUDA_USE_NVIDIA_BINDING`` to ``1`` prior to the import of Numba. Once
-Numba has been imported, the selected binding cannot be changed.
+Numba-CUDA uses the official `NVIDIA CUDA Python bindings
+<https://nvidia.github.io/cuda-python/>`_ for all CUDA Driver interactions.
+Numba-CUDA previously provided its own internal ctypes-based bindings; the
+public APIs exposing those bindings are kept for compatibility, but if you
+need to interact directly with the CUDA Driver or other CUDA libraries we
+recommend using the `cuda-python <https://nvidia.github.io/cuda-python/>`_
+package directly.
 
 
 Per-Thread Default Streams
 --------------------------
 
 Responsibility for handling Per-Thread Default Streams (PTDS) is delegated to
-the NVIDIA bindings when they are in use. To use PTDS with the NVIDIA bindings,
-set the environment variable ``CUDA_PYTHON_CUDA_PER_THREAD_DEFAULT_STREAM`` to
-``1`` instead of Numba's environmnent variable
-:envvar:`NUMBA_CUDA_PER_THREAD_DEFAULT_STREAM`.
+the NVIDIA bindings. To use PTDS, set the environment variable
+``CUDA_PYTHON_CUDA_PER_THREAD_DEFAULT_STREAM`` to ``1`` instead of Numba's
+environment variable :envvar:`NUMBA_CUDA_PER_THREAD_DEFAULT_STREAM`.
 
 .. seealso::
 
@@ -35,13 +32,5 @@ set the environment variable ``CUDA_PYTHON_CUDA_PER_THREAD_DEFAULT_STREAM`` to
 Roadmap
 -------
 
-In Numba 0.56, the NVIDIA Bindings will be used by default, if they are
-installed.
-
-In future versions of Numba:
-
-- The internal bindings will be deprecated.
-- The internal bindings will be removed.
-
-At present, no specific release is planned for the deprecation or removal of
-the internal bindings.
+The ctypes-based internal bindings have been removed in favor of the NVIDIA
+bindings. Future work focuses on expanding usage of ``cuda.core`` APIs.

docs/source/reference/envvars.rst

@@ -61,14 +61,12 @@ Configuration
 CUDA Bindings
 -------------
 
-Numba supports interacting with the CUDA Driver API via either the `NVIDIA CUDA
-Python bindings <https://nvidia.github.io/cuda-python/>`_ or its own ctypes-based
-bindings. Functionality is equivalent between the two binding choices. The
-NVIDIA bindings are the default, and the ctypes bindings are now deprecated.
-
-If you do not want to use the NVIDIA bindings, the (deprecated) ctypes bindings
-can be enabled by setting the environment variable
-:envvar:`NUMBA_CUDA_USE_NVIDIA_BINDING` to ``"0"``.
+Numba-CUDA uses the `NVIDIA CUDA Python bindings <https://nvidia.github.io/cuda-python/>`_
+for interacting with the CUDA Driver API. Numba-CUDA previously provided its own
+internal ctypes-based bindings; the public APIs exposing those bindings are kept
+for compatibility, but if you need to interact directly with the CUDA Driver or
+other CUDA libraries we recommend using the `cuda-python <https://nvidia.github.io/cuda-python/cuda-pathfinder/latest/generated/cuda.pathfinder.load_nvidia_dynamic_lib.html>`_
+package directly.
 
 
 .. _cudatoolkit-lookup:
@@ -79,22 +77,8 @@ CUDA Driver and Toolkit search paths
 Default behavior
 ~~~~~~~~~~~~~~~~
 
-When using the NVIDIA bindings, searches for the CUDA driver and toolkit
-libraries use its `built-in path-finding logic <https://github.com/NVIDIA/cuda-python/tree/main/cuda_bindings/cuda/bindings/_path_finder>`_.
-
-Ctypes bindings (deprecated) behavior
-~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
-
-When using the ctypes bindings, Numba searches for a CUDA toolkit installation
-in the following order:
-
-1. Conda-installed CUDA Toolkit packages
-2. Pip-installed CUDA Toolkit packages
-3. The environment variable ``CUDA_HOME``, which points to the directory of the
-   installed CUDA toolkit (i.e. ``/home/user/cuda-12``)
-4. System-wide installation at exactly ``/usr/local/cuda`` on Linux platforms.
-   Versioned installation paths (i.e. ``/usr/local/cuda-12.0``) are intentionally
-   ignored. Users can use ``CUDA_HOME`` to select specific versions.
+Searches for the CUDA driver and toolkit libraries use the NVIDIA bindings'
+`built-in path-finding logic <https://github.com/NVIDIA/cuda-python/tree/main/cuda_bindings/cuda/bindings/_path_finder>`_.
 
 In addition to the CUDA toolkit libraries, which can be installed by conda into
 an environment or installed system-wide by the `CUDA SDK installer

docs/source/user/bindings.rst

@@ -8,65 +8,15 @@
 import sys
 
 
-# Enable pynvjitlink based on the following precedence:
-# 1. Config setting "CUDA_ENABLE_PYNVJITLINK" (highest priority)
-# 2. Environment variable "NUMBA_CUDA_ENABLE_PYNVJITLINK"
-# 3. Auto-detection of pynvjitlink module (lowest priority)
-
-pynvjitlink_auto_enabled = False
-
-if getattr(config, "CUDA_ENABLE_PYNVJITLINK", None) is None:
-    if (
-        _pynvjitlink_enabled_in_env := _readenv(
-            "NUMBA_CUDA_ENABLE_PYNVJITLINK", bool, None
-        )
-    ) is not None:
-        config.CUDA_ENABLE_PYNVJITLINK = _pynvjitlink_enabled_in_env
-    else:
-        pynvjitlink_auto_enabled = (
-            importlib.util.find_spec("pynvjitlink") is not None
-        )
-        config.CUDA_ENABLE_PYNVJITLINK = pynvjitlink_auto_enabled
-
-# Upstream numba sets CUDA_USE_NVIDIA_BINDING to 0 by default, so it always
-# exists. Override, but not if explicitly set to 0 in the envioronment.
-_nvidia_binding_enabled_in_env = _readenv(
-    "NUMBA_CUDA_USE_NVIDIA_BINDING", bool, None
-)
-if _nvidia_binding_enabled_in_env is False:
-    USE_NV_BINDING = False
-else:
-    USE_NV_BINDING = True
-    config.CUDA_USE_NVIDIA_BINDING = USE_NV_BINDING
-if config.CUDA_USE_NVIDIA_BINDING:
-    if not (
-        importlib.util.find_spec("cuda")
-        and importlib.util.find_spec("cuda.bindings")
-    ):
-        raise ImportError(
-            "CUDA bindings not found. Please pip install the "
-            "cuda-bindings package. Alternatively, install "
-            "numba-cuda[cuXY], where XY is the required CUDA "
-            "version, to install the binding automatically. "
-            "If no CUDA bindings are desired, set the env var "
-            "NUMBA_CUDA_USE_NVIDIA_BINDING=0 to enable ctypes "
-            "bindings."
-        )
-
-if config.CUDA_ENABLE_PYNVJITLINK:
-    if USE_NV_BINDING and not pynvjitlink_auto_enabled:
-        warnings.warn(
-            "Explicitly enabling pynvjitlink is no longer necessary. "
-            "NVIDIA bindings are enabled. cuda.core will be used "
-            "in place of pynvjitlink."
-        )
-    elif pynvjitlink_auto_enabled:
-        # Ignore the fact that pynvjitlink is enabled, because that was an
-        # automatic decision based on discovering pynvjitlink was present; the
-        # user didn't ask for it
-        pass
-    else:
-        raise RuntimeError("nvJitLink requires the NVIDIA CUDA bindings. ")
+# Require NVIDIA CUDA bindings at import time
+if not (
+    importlib.util.find_spec("cuda")
+    and importlib.util.find_spec("cuda.bindings")
+):
+    raise ImportError(
+        "NVIDIA CUDA Python bindings not found. Install the 'cuda' package "
+        "(e.g. pip install nvidia-cuda-python or numba-cuda[cuXY])."
+    )
 
 if config.ENABLE_CUDASIM:
     from .simulator_init import *

docs/source/user/installation.rst

@@ -8,7 +8,6 @@
 import contextlib
 import os
 
-from numba.cuda.cudadrv import drvapi
 import numpy as np
 
 from .cudadrv import devicearray, devices, driver
@@ -48,10 +47,7 @@ def from_cuda_array_interface(desc, owner=None, sync=True):
     )
     size = driver.memory_size_from_info(shape, strides, dtype.itemsize)
 
-    if config.CUDA_USE_NVIDIA_BINDING:
-        cudevptr_class = driver.binding.CUdeviceptr
-    else:
-        cudevptr_class = drvapi.cu_device_ptr
+    cudevptr_class = driver.binding.CUdeviceptr
     devptr = cudevptr_class(desc["data"][0])
     data = driver.MemoryPointer(
         current_context(), devptr, size=size, owner=owner

numba_cuda/numba/cuda/__init__.py

@@ -26,10 +26,7 @@ def run_nvdisasm(cubin, flags):
     try:
         fd, fname = tempfile.mkstemp()
         with open(fname, "wb") as f:
-            if config.CUDA_USE_NVIDIA_BINDING:
-                f.write(cubin.code)
-            else:
-                f.write(cubin)
+            f.write(cubin.code)
 
         try:
             cp = subprocess.run(