Repository: numba/numba

Issue:
Retrieve The Oldest and Newest Supported CUDA Versions
## Feature request

The STUMPY Python package depends on `numba` and, when `cuda` is available, it imports submodules that rely on cuda. It checks this by doing something like:

```
from numba import cuda

if cuda.is_available():  # Only import if CUDA is available
    from stumpy import gpu_stump
```

However, `numba` does not currently support CUDA 11.2 and if CUDA 11.2 is installed then the above `if` condition still returns `True`. It would be nice to have function that checks if the available CUDA version is compatible with `numba` like:

```
if cuda.is_available() and cuda.version_is_supported():
    from stumpy import gpu_stump
```

Or, as discussed with @stuartarchibald, maybe it makes more sense to add max/min supported versions enumerated and then make that part of the `cuda.is_available()` check. 

Discussion:

Patch:
diff --git a/docs/source/cuda-reference/host.rst b/docs/source/cuda-reference/host.rst
index d69b5de4e..9be7e7773 100644
--- a/docs/source/cuda-reference/host.rst
+++ b/docs/source/cuda-reference/host.rst
@@ -213,4 +213,9 @@ accessed through ``cuda.runtime``, which is an instance of the
 :class:`numba.cuda.cudadrv.runtime.Runtime` class:
 
 .. autoclass:: numba.cuda.cudadrv.runtime.Runtime
-   :members: get_version
+   :members: get_version, is_supported_version, supported_versions
+
+Whether the current runtime is officially supported and tested with the current
+version of Numba can also be queried:
+
+.. autofunction:: numba.cuda.is_supported_version
diff --git a/docs/source/cuda/overview.rst b/docs/source/cuda/overview.rst
index f72341840..d4575487c 100644
--- a/docs/source/cuda/overview.rst
+++ b/docs/source/cuda/overview.rst
@@ -45,7 +45,11 @@ Software
 
 Numba aims to support CUDA Toolkit versions released within the last 3 years. At
 the present time, you will need the CUDA toolkit version 9.0 or later installed.
-If you are using Conda, just type::
+
+CUDA is supported on 64-bit Linux and Windows. 32-bit platforms, and macOS are
+unsupported.
+
+If you are using Conda, you can install the CUDA toolkit with::
 
    $ conda install cudatoolkit
 
diff --git a/docs/source/user/installing.rst b/docs/source/user/installing.rst
index 85c60a2f9..4b7bd843e 100644
--- a/docs/source/user/installing.rst
+++ b/docs/source/user/installing.rst
@@ -14,7 +14,7 @@ Our supported platforms are:
 * Windows 7 and later (32-bit and 64-bit)
 * OS X 10.9 and later (64-bit)
 * \*BSD (unofficial support only)
-* NVIDIA GPUs of compute capability 2.0 and later
+* NVIDIA GPUs of compute capability 3.0 and later
 * AMD ROC dGPUs (linux only and not for AMD Carrizo or Kaveri APU)
 * ARMv7 (32-bit little-endian, such as Raspberry Pi 2 and 3)
 * ARMv8 (64-bit little-endian, such as the NVIDIA Jetson)
diff --git a/numba/cuda/cudadrv/runtime.py b/numba/cuda/cudadrv/runtime.py
index b5c2139bb..a72f7dc69 100644
--- a/numba/cuda/cudadrv/runtime.py
+++ b/numba/cuda/cudadrv/runtime.py
@@ -7,6 +7,7 @@ really used in Numba except for querying the Runtime version.
 
 import ctypes
 import functools
+import sys
 
 from numba.core import config
 from numba.cuda.cudadrv.driver import ERROR_MAP, make_logger
@@ -113,5 +114,23 @@ class Runtime:
         minor = (rtver.value - (major * 1000)) // 10
         return (major, minor)
 
+    def is_supported_version(self):
+        """
+        Returns True if the CUDA Runtime is a supported version.
+        """
+
+        return self.get_version() in self.supported_versions
+
+    @property
+    def supported_versions(self):
+        """A tuple of all supported CUDA toolkit versions. Versions are given in
+        the form ``(major_version, minor_version)``."""
+        if sys.platform not in ('linux', 'win32') or config.MACHINE_BITS != 64:
+            # Only 64-bit Linux and Windows are supported
+            return ()
+        return ((9, 0), (9, 1), (9, 2),
+                (10, 0), (10, 1), (10, 2),
+                (11, 0), (11, 1), (11, 2))
+
 
 runtime = Runtime()
diff --git a/numba/cuda/device_init.py b/numba/cuda/device_init.py
index 7566f40fc..be5f46b8c 100644
--- a/numba/cuda/device_init.py
+++ b/numba/cuda/device_init.py
@@ -51,6 +51,23 @@ def is_available():
     return driver_is_available and nvvm.is_available()
 
 
+def is_supported_version():
+    """Returns True if the CUDA Runtime is a supported version.
+
+    Unsupported versions (e.g. newer versions than those known to Numba)
+    may still work; this function provides a facility to check whether the
+    current Numba version is tested and known to work with the current
+    runtime version. If the current version is unsupported, the caller can
+    decide how to act. Options include:
+
+    - Continuing silently,
+    - Emitting a warning,
+    - Generating an error or otherwise preventing the use of CUDA.
+    """
+
+    return runtime.is_supported_version()
+
+
 def cuda_error():
     """Returns None or an exception if the CUDA driver fails to initialize.
     """
diff --git a/numba/cuda/simulator/cudadrv/runtime.py b/numba/cuda/simulator/cudadrv/runtime.py
index c88ae2aad..308d19e76 100644
--- a/numba/cuda/simulator/cudadrv/runtime.py
+++ b/numba/cuda/simulator/cudadrv/runtime.py
@@ -8,5 +8,12 @@ class FakeRuntime(object):
     def get_version(self):
         return (-1, -1)
 
+    def is_supported_version(self):
+        return True
+
+    @property
+    def supported_versions(self):
+        return (-1, -1),
+
 
 runtime = FakeRuntime()


Test patch:
diff --git a/numba/cuda/tests/cudadrv/test_runtime.py b/numba/cuda/tests/cudadrv/test_runtime.py
index 967a4b625..ad4a85cde 100644
--- a/numba/cuda/tests/cudadrv/test_runtime.py
+++ b/numba/cuda/tests/cudadrv/test_runtime.py
@@ -2,7 +2,8 @@ import multiprocessing
 import os
 from numba.core import config
 from numba.cuda.cudadrv.runtime import runtime
-from numba.cuda.testing import unittest, SerialMixin
+from numba.cuda.testing import unittest, SerialMixin, skip_on_cudasim
+from unittest.mock import patch
 
 
 def set_visible_devices_and_check(q):
@@ -17,14 +18,33 @@ def set_visible_devices_and_check(q):
         q.put(-1)
 
 
+if config.ENABLE_CUDASIM:
+    SUPPORTED_VERSIONS = (-1, -1),
+else:
+    SUPPORTED_VERSIONS = ((9, 0), (9, 1), (9, 2),
+                          (10, 0), (10, 1), (10, 2),
+                          (11, 0), (11, 1), (11, 2))
+
+
 class TestRuntime(unittest.TestCase):
     def test_get_version(self):
-        if config.ENABLE_CUDASIM:
-            supported_versions = (-1, -1),
-        else:
-            supported_versions = ((9, 0), (9, 1), (9, 2), (10, 0),
-                                  (10, 1), (10, 2), (11, 0), (11, 1), (11, 2))
-        self.assertIn(runtime.get_version(), supported_versions)
+        self.assertIn(runtime.get_version(), SUPPORTED_VERSIONS)
+
+    def test_is_supported_version_true(self):
+        for v in SUPPORTED_VERSIONS:
+            with patch.object(runtime, 'get_version', return_value=v):
+                self.assertTrue(runtime.is_supported_version())
+
+    @skip_on_cudasim('The simulator always simulates a supported runtime')
+    def test_is_supported_version_false(self):
+        # Check with an old unsupported version and some potential future
+        # versions
+        for v in ((8, 0), (11, 3), (12, 0)):
+            with patch.object(runtime, 'get_version', return_value=v):
+                self.assertFalse(runtime.is_supported_version())
+
+    def test_supported_versions(self):
+        self.assertEqual(SUPPORTED_VERSIONS, runtime.supported_versions)
 
 
 class TestVisibleDevices(unittest.TestCase, SerialMixin):
diff --git a/numba/cuda/tests/cudapy/test_constmem.py b/numba/cuda/tests/cudapy/test_constmem.py
index f32700aa0..89cab48df 100644
--- a/numba/cuda/tests/cudapy/test_constmem.py
+++ b/numba/cuda/tests/cudapy/test_constmem.py
@@ -138,7 +138,7 @@ class TestCudaConstantMemory(CUDATestCase):
             # and > 11.1 use a vector of 2x f32. The root cause of these
             # codegen differences is not known, but must be accounted for in
             # this test.
-            if cuda.runtime.get_version() in ((8, 0), (9, 0), (9, 1), (11, 2)):
+            if cuda.runtime.get_version() in ((9, 0), (9, 1), (11, 2)):
                 complex_load = 'ld.const.v2.f32'
                 description = 'Load the complex as a vector of 2x f32'
             else: