sklam · Feb 25, 2023 · Jun 21, 2022 · Jul 21, 2022 · Jul 21, 2022 · Sep 1, 2022
diff --git a/buildscripts/gpuci/axis.yaml b/buildscripts/gpuci/axis.yaml
@@ -9,6 +9,7 @@ CUDA_TOOLKIT_VER:
 - "11.1"
 - "11.2"
 - "11.5"
+- "11.8"
 
 LINUX_VER:
 - ubuntu18.04

diff --git a/buildscripts/gpuci/build.sh b/buildscripts/gpuci/build.sh
@@ -24,10 +24,18 @@ else
   export NUMBA_CUDA_USE_NVIDIA_BINDING=0;
 fi;
 
+# Test with Minor Version Compatibility on CUDA 11.8
+if [ $CUDA_TOOLKIT_VER == "11.8" ]
+then
+  export NUMBA_CUDA_ENABLE_MINOR_VERSION_COMPATIBILITY=1;
+else
+  export NUMBA_CUDA_ENABLE_MINOR_VERSION_COMPATIBILITY=0;
+fi;
+
 # Test with different NumPy versions with each toolkit (it's not worth testing
 # the Cartesian product of versions here, we just need to test with different
 # CUDA and NumPy versions).
-declare -A CTK_NUMPY_VMAP=( ["11.0"]="1.19" ["11.1"]="1.21" ["11.2"]="1.22" ["11.5"]="1.23")
+declare -A CTK_NUMPY_VMAP=( ["11.0"]="1.19" ["11.1"]="1.20" ["11.2"]="1.21" ["11.5"]="1.22" ["11.8"]="1.23")
 NUMPY_VER="${CTK_NUMPY_VMAP[$CUDA_TOOLKIT_VER]}"
 
 
@@ -46,6 +54,8 @@ gpuci_logger "Create testing env"
 gpuci_mamba_retry create -n numba_ci -y \
                   "python=${PYTHON_VER}" \
                   "cudatoolkit=${CUDA_TOOLKIT_VER}" \
+                  "rapidsai::cubinlinker" \
+                  "conda-forge::ptxcompiler" \
                   "numba/label/dev::llvmlite" \
                   "numpy=${NUMPY_VER}" \
                   "scipy" \

diff --git a/docs/source/cuda/index.rst b/docs/source/cuda/index.rst
@@ -26,4 +26,5 @@ Numba for CUDA GPUs
    bindings.rst
    cuda_ffi.rst
    caching.rst
+   minor_version_compatibility.rst
    faq.rst
diff --git a/docs/source/cuda/minor_version_compatibility.rst b/docs/source/cuda/minor_version_compatibility.rst
@@ -0,0 +1,64 @@
+.. _minor-version-compatibility:
+
+CUDA Minor Version Compatiblity
+===============================
+
+CUDA `Minor Version Compatibility
+<https://docs.nvidia.com/deploy/cuda-compatibility/index.html#minor-version-compatibility>`_
+(MVC) enables the use of a newer CUDA toolkit version than the CUDA version
+supported by the driver, provided that the toolkit and driver both have the same
+major version. For example, use of CUDA toolkit 11.5 with CUDA driver 450 (CUDA
+version 11.0) is supported through MVC.
+
+Numba supports MVC for CUDA 11 on Linux using the external ``cubinlinker`` and
+``ptxcompiler`` packages, subject to the following limitations:
+
+- Linking of archives is unsupported.
+- Cooperative Groups are unsupported, because they require an archive to be
+  linked.
+
+MVC is not supported on Windows.
+
+
+Enabling MVC Support
+--------------------
+
+To use MVC support, the ``cubinlinker`` and ``ptxcompiler`` compiler packages
+must be installed from the appropriate channels. To install using conda, use:
+
+.. code:: bash
+
+   conda install rapidsai::cubinlinker conda-forge::ptxcompiler
+
+To install with pip, use the NVIDIA package index:
+
+.. code:: bash
+
+   pip install ptxcompiler-cu11 cubinlinker-cu11 --extra-index-url=https://pypi.ngc.nvidia.com
+
+MVC support is enabled by setting the environment variable:
+
+.. code:: bash
+
+   export NUMBA_CUDA_ENABLE_MINOR_VERSION_COMPATIBILITY=1
+
+
+or by setting a configuration variable prior to using any CUDA functionality in
+Numba:
+
+.. code:: python
+
+   from numba import config
+   config.CUDA_ENABLE_MINOR_VERSION_COMPATIBILITY = True
+
+
+References
+----------
+
+Further information about Minor Version Compatibility may be found in:
+
+- The `CUDA Compatibility Guide
+  <https://docs.nvidia.com/deploy/cuda-compatibility/index.html>`_.
+- The `README for ptxcompiler
+  <https://github.com/rapidsai/ptxcompiler/blob/main/README.md>`_.
+
diff --git a/docs/source/cuda/overview.rst b/docs/source/cuda/overview.rst
@@ -55,8 +55,8 @@ Software
 --------
 
 Numba aims to support CUDA Toolkit versions released within the last 3 years.
-An NVIDIA driver sufficient for the toolkit version is also required.
-Presently:
+An NVIDIA driver sufficient for the toolkit version is also required (see also
+:ref:`minor-version-compatibility`). Presently:
 
 * 11.0 is the minimum required toolkit version.
 * 11.2 or later is recommended, as it uses an NVVM version based on LLVM 7 (as

diff --git a/docs/source/user/installing.rst b/docs/source/user/installing.rst
@@ -233,6 +233,9 @@ vary with target operating system and hardware. The following lists them all
     :ref:`runtime type-checking <type_anno_check>`.
   * ``cuda-python`` - The NVIDIA CUDA Python bindings. See :ref:`cuda-bindings`.
     Numba requires Version 11.6 or greater.
+  * ``cubinlinker`` and ``ptxcompiler`` to support
+    :ref:`minor-version-compatibility`.
+
 
 * To build the documentation:
 

diff --git a/numba/core/config.py b/numba/core/config.py
@@ -419,6 +419,9 @@ def avx_default():
         CUDA_PER_THREAD_DEFAULT_STREAM = _readenv(
             "NUMBA_CUDA_PER_THREAD_DEFAULT_STREAM", int, 0)
 
+        CUDA_ENABLE_MINOR_VERSION_COMPATIBILITY = _readenv(
+            "NUMBA_CUDA_ENABLE_MINOR_VERSION_COMPATIBILITY", int, 0)
+
         # Location of the CUDA include files
         if IS_WIN32:
             cuda_path = os.environ.get('CUDA_PATH')

@@ -6,8 +6,6 @@
 from numba.core.errors import NumbaInvalidConfigWarning
 from .cudadrv import devices, driver, nvvm, runtime
 
-import ctypes
-import numpy as np
 import os
 import subprocess
 import tempfile
@@ -181,11 +179,7 @@ def get_cubin(self, cc=None):
         for path in self._linking_files:
             linker.add_file_guess_ext(path)
 
-        cubin_buf, size = linker.complete()
-
-        # We take a copy of the cubin because it's owned by the linker
-        cubin_ptr = ctypes.cast(cubin_buf, ctypes.POINTER(ctypes.c_char))
-        cubin = bytes(np.ctypeslib.as_array(cubin_ptr, shape=(size,)))
+        cubin = linker.complete()
         self._cubin_cache[cc] = cubin
         self._linkerinfo_cache[cc] = linker.info_log
 

@@ -20,6 +20,7 @@
 import logging
 import threading
 import asyncio
+import pathlib
 from itertools import product
 from abc import ABCMeta, abstractmethod
 from ctypes import (c_int, byref, c_size_t, c_char, c_char_p, addressof,
@@ -36,6 +37,14 @@
 from .drvapi import cu_occupancy_b2d_size, cu_stream_callback_pyobj, cu_uuid
 from numba.cuda.cudadrv import enums, drvapi, _extras
 
+if config.CUDA_ENABLE_MINOR_VERSION_COMPATIBILITY:
+    try:
+        from ptxcompiler import compile_ptx
+        from cubinlinker import CubinLinker, CubinLinkerError
+    except ImportError as ie:
+        msg = ("Minor version compatibility requires ptxcompiler and "
+               "cubinlinker packages to be available")
+        raise ImportError(msg) from ie
 
 USE_NV_BINDING = config.CUDA_USE_NVIDIA_BINDING
 
@@ -2584,7 +2593,9 @@ class Linker(metaclass=ABCMeta):
 
     @classmethod
     def new(cls, max_registers=0, lineinfo=False, cc=None):
-        if USE_NV_BINDING:
+        if config.CUDA_ENABLE_MINOR_VERSION_COMPATIBILITY:
+            return MVCLinker(max_registers, lineinfo, cc)
+        elif USE_NV_BINDING:
             return CudaPythonLinker(max_registers, lineinfo, cc)
         else:
             return CtypesLinker(max_registers, lineinfo, cc)
@@ -2644,6 +2655,85 @@ def complete(self):
         """
 
 
+class MVCLinker(Linker):
+    """
+    Linker supporting Minor Version Compatibility, backed by the cubinlinker
+    package.
+    """
+    def __init__(self, max_registers=None, lineinfo=False, cc=None):
+        if cc is None:
+            raise RuntimeError("MVCLinker requires Compute Capability to be "
+                               "specified, but cc is None")
+
+        arch = f"sm_{cc[0] * 10 + cc[1]}"
+        ptx_compile_opts = ['--gpu-name', arch, '-c']
+        if max_registers:
+            arg = f"--maxrregcount={max_registers}"
+            ptx_compile_opts.append(arg)
+        if lineinfo:
+            ptx_compile_opts.append('--generate-line-info')
+        self.ptx_compile_options = tuple(ptx_compile_opts)
+
+        self._linker = CubinLinker(f"--arch={arch}")
+
+    @property
+    def info_log(self):
+        return self._linker.info_log
+
+    @property
+    def error_log(self):
+        return self._linker.error_log
+
+    def add_ptx(self, ptx, name='<cudapy-ptx>'):
+        compile_result = compile_ptx(ptx.decode(), self.ptx_compile_options)
+        try:
+            self._linker.add_cubin(compile_result.compiled_program, name)
+        except CubinLinkerError as e:
+            raise LinkerError from e
+
+    def add_file(self, path, kind):
+        try:
+            with open(path, 'rb') as f:
+                data = f.read()
+        except FileNotFoundError:
+            raise LinkerError(f'{path} not found')
+
+        name = pathlib.Path(path).name
+        if kind == FILE_EXTENSION_MAP['cubin']:
+            fn = self._linker.add_cubin
+        elif kind == FILE_EXTENSION_MAP['fatbin']:
+            fn = self._linker.add_fatbin
+        elif kind == FILE_EXTENSION_MAP['a']:
+            raise LinkerError(f"Don't know how to link {kind}")
+        elif kind == FILE_EXTENSION_MAP['ptx']:
+            return self.add_ptx(data, name)
+        else:
+            raise LinkerError(f"Don't know how to link {kind}")
+
+        try:
+            fn(data, name)
+        except CubinLinkerError as e:
+            raise LinkerError from e
+
+    def add_cu(self, cu, name):
+        program = NvrtcProgram(cu, name)
+
+        if config.DUMP_ASSEMBLY:
+            print(("ASSEMBLY %s" % name).center(80, '-'))
+            print(program.ptx.decode())
+            print('=' * 80)
+
+        # Link the program's PTX using the normal linker mechanism
+        ptx_name = os.path.splitext(name)[0] + ".ptx"
+        self.add_ptx(program.ptx.rstrip(b'\x00'), ptx_name)
+
+    def complete(self):
+        try:
+            return self._linker.complete()
+        except CubinLinkerError as e:
+            raise LinkerError from e
+
+
 class CtypesLinker(Linker):
     """
     Links for current device if no CC given
@@ -2725,18 +2815,21 @@ def add_cu(self, path, name):
                                   "with the ctypes binding. ")
 
     def complete(self):
-        cubin = c_void_p(0)
+        cubin_buf = c_void_p(0)
         size = c_size_t(0)
 
         try:
-            driver.cuLinkComplete(self.handle, byref(cubin), byref(size))
+            driver.cuLinkComplete(self.handle, byref(cubin_buf), byref(size))
         except CudaAPIError as e:
             raise LinkerError("%s\n%s" % (e, self.error_log))
 
         size = size.value
         assert size > 0, 'linker returned a zero sized cubin'
         del self._keep_alive[:]
-        return cubin, size
+
+        # We return a copy of the cubin because it's owned by the linker
+        cubin_ptr = ctypes.cast(cubin_buf, ctypes.POINTER(ctypes.c_char))
+        return bytes(np.ctypeslib.as_array(cubin_ptr, shape=(size,)))
 
 
 class NvrtcProgram:
@@ -2911,13 +3004,15 @@ def add_file(self, path, kind):
 
     def complete(self):
         try:
-            cubin, size = driver.cuLinkComplete(self.handle)
+            cubin_buf, size = driver.cuLinkComplete(self.handle)
         except CudaAPIError as e:
             raise LinkerError("%s\n%s" % (e, self.error_log))
 
         assert size > 0, 'linker returned a zero sized cubin'
         del self._keep_alive[:]
-        return cubin, size
+        # We return a copy of the cubin because it's owned by the linker
+        cubin_ptr = ctypes.cast(cubin_buf, ctypes.POINTER(ctypes.c_char))
+        return bytes(np.ctypeslib.as_array(cubin_ptr, shape=(size,)))
 
 
 # -----------------------------------------------------------------------------

@@ -109,6 +109,12 @@ def skip_if_cuda_includes_missing(fn):
     return unittest.skipUnless(cuda_h_file, reason)(fn)
 
 
+def skip_if_mvc_enabled(reason):
+    """Skip a test if Minor Version Compatibility is enabled"""
+    return unittest.skipIf(config.CUDA_ENABLE_MINOR_VERSION_COMPATIBILITY,
+                           reason)
+
+
 def cc_X_or_above(major, minor):
     if not config.ENABLE_CUDASIM:
         cc = devices.get_context().device.compute_capability

@@ -97,7 +97,7 @@ class TestLinker(CUDATestCase):
     def test_linker_basic(self):
         '''Simply go through the constructor and destructor
         '''
-        linker = Linker.new()
+        linker = Linker.new(cc=(5, 3))
         del linker
 
     def _test_linking(self, eager):

@@ -4,7 +4,8 @@
 
 from numba import config, cuda, int32
 from numba.cuda.testing import (unittest, CUDATestCase, skip_on_cudasim,
-                                skip_unless_cc_60, skip_if_cudadevrt_missing)
+                                skip_unless_cc_60, skip_if_cudadevrt_missing,
+                                skip_if_mvc_enabled)
 
 
 @cuda.jit
@@ -46,6 +47,7 @@ def sequential_rows(M):
 
 
 @skip_if_cudadevrt_missing
+@skip_if_mvc_enabled('CG not supported with MVC')
 class TestCudaCooperativeGroups(CUDATestCase):
     @skip_unless_cc_60
     def test_this_grid(self):

@@ -3,11 +3,13 @@
 
 import unittest
 from numba.cuda.testing import (CUDATestCase, skip_on_cudasim,
-                                skip_if_cudadevrt_missing, skip_unless_cc_60)
+                                skip_if_cudadevrt_missing, skip_unless_cc_60,
+                                skip_if_mvc_enabled)
 
 
 @skip_if_cudadevrt_missing
 @skip_unless_cc_60
+@skip_if_mvc_enabled('CG not supported with MVC')
 @skip_on_cudasim("cudasim doesn't support cuda import at non-top-level")
 class TestCooperativeGroups(CUDATestCase):
     def test_ex_grid_sync(self):

@@ -1,12 +1,14 @@
 import unittest
 
 from numba.cuda.testing import (CUDATestCase, skip_if_cudadevrt_missing,
-                                skip_on_cudasim, skip_unless_cc_60)
+                                skip_on_cudasim, skip_unless_cc_60,
+                                skip_if_mvc_enabled)
 from numba.tests.support import captured_stdout
 
 
 @skip_if_cudadevrt_missing
 @skip_unless_cc_60
+@skip_if_mvc_enabled('CG not supported with MVC')
 @skip_on_cudasim("cudasim doesn't support cuda import at non-top-level")
 class TestLaplace(CUDATestCase):
     """
-Original file line number
+Diff line change
@@ Expand Up / @@ -9,6 +9,7 @@ CUDA_TOOLKIT_VER: @@
     - "11.1"
     - "11.2"
     - "11.5"
+    - "11.8"
     LINUX_VER:
     - ubuntu18.04
@@ Expand Down @@