Repository: numba/numba
Pull request: CUDA: Add a function to query whether the runtime version is supported.
`numba.cuda.is_supported_version()` returns `True` if the CUDA Runtime is a supported version (as suggested in #6645). The list of supported versions matches those with which the current version of Numba is tested.

Unsupported versions (e.g. newer versions than those known to Numba) may still work. 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.


Notes:

- Based on #6661, which adds support for CUDA 11.2
- Also removes a reference to an unsupported version (8.0) in `test_constmem`.
- Fixes #6645.

Problem statement:
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.

Hints from issue discussion:
I would also agree with the approach of combining the version check with the driver check in `cuda.is_available()` to keep things simple for downstream libraries.
Thanks for writing this up @seanlaw, original discussion is here: https://gitter.im/numba/numba?at=6001c07114cec811ec9e267a, as noted in that, CUDA `version++` has not in my memory broken Numba support before which is why sadly these support querying features don't exist. I think this needs to go into 0.53.0 to make sure there's a way to do this for the future.
CC @gmarkall @testhound 
This is addressed in #6666 - After some discussion with @stuartarchibald, I didn't merge the supported version check into `is_available()`, because 

- One might want to proceed even if the version isn't a known-good version (or may want to just emit a warning).
- These kinds of breakages are rare - it's only happened once so far, and is partly due to Numba using a deprecated interface to NVVM for a long time after it was deprecated.
Thanks for the update and for working on this, @gmarkall! I think having a separate `numba.cuda.is_supported_version()` function is perfectly reasonable. Is it possible to add a parameter/argument to the function call to retrieve the min/max compatible versions (or maybe another function)? This way, I can emit a warning to explicitly say "the cuda version that you are using is X and is newer than the latest supported version Y". But, right now, the function appears to only return a `bool`
@seanlaw I've added a property to get the list of supported versions, from which the max and min supported versions can be derived - I thought about max / min, but just in case we end up with an unsupported version in the middle of the range for some reason, the API won't need changing  in future.
Fantastic! Thanks, @gmarkall - please feel free to close this issue

Commit message: Don't advertise CUDA version support on unsupported platforms

Patch to review:
diff --git a/docs/source/cuda-reference/host.rst b/docs/source/cuda-reference/host.rst
index d69b5de4e48..9be7e7773aa 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 f7234184039..d4575487cc2 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 85c60a2f9bd..4b7bd843e3c 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 b5c2139bb3c..fa63cab77bc 100644
--- a/numba/cuda/cudadrv/runtime.py
+++ b/numba/cuda/cudadrv/runtime.py
@@ -7,6 +7,7 @@
 
 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 @@ def get_version(self):
         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 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 7566f40fc55..b6463074b47 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 c88ae2aadd1..308d19e7683 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()
diff --git a/numba/cuda/tests/cudadrv/test_runtime.py b/numba/cuda/tests/cudadrv/test_runtime.py
index 967a4b62564..ad4a85cdee2 100644
--- a/numba/cuda/tests/cudadrv/test_runtime.py
+++ b/numba/cuda/tests/cudadrv/test_runtime.py
@@ -2,7 +2,8 @@
 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 f32700aa001..89cab48dfef 100644
--- a/numba/cuda/tests/cudapy/test_constmem.py
+++ b/numba/cuda/tests/cudapy/test_constmem.py
@@ -138,7 +138,7 @@ def test_const_array_3d(self):
             # 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:


Code review comments:
diff_hunk: @@ -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.
path: numba/cuda/device_init.py
text: @user1:
```suggestion
    """Returns True if the CUDA Runtime is a supported version.
```