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
2 changes: 1 addition & 1 deletion buildscripts/gpuci/axis.yaml
Original file line number Diff line number Diff line change
Expand Up @@ -5,7 +5,7 @@ CUDA_VER:
- "11.2"

CUDA_TOOLKIT_VER:
- "10.2"
- "11.0"
- "11.1"
- "11.2"
- "11.5"
Expand Down
5 changes: 1 addition & 4 deletions docs/source/cuda/cooperative_groups.rst
Original file line number Diff line number Diff line change
Expand Up @@ -11,10 +11,7 @@ synchronization, along with cooperative kernel launches.
Cooperative groups are supported on Linux, and Windows for devices in `TCC
mode
<https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#tesla-compute-cluster-mode-for-windows>`_.
Cooperative Groups also require the CUDA Device Runtime library, ``cudadevrt``,
to be available - for conda default channel-installed CUDA toolkit packages, it
is only available in versions 10.2 onwards. System-installed toolkits (e.g. from
NVIDIA distribution packages or runfiles) all include ``cudadevrt``.


Using Grid Groups
-----------------
Expand Down
2 changes: 1 addition & 1 deletion docs/source/cuda/overview.rst
Original file line number Diff line number Diff line change
Expand Up @@ -58,7 +58,7 @@ 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:

* 10.2 is the minimum required toolkit version.
* 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
opposed to 3.4 in earlier releases).

Expand Down
9 changes: 5 additions & 4 deletions docs/source/reference/deprecation.rst
Original file line number Diff line number Diff line change
Expand Up @@ -289,10 +289,10 @@ Schedule
- In Numba 0.56: ``add_user_function()`` was removed.


Deprecation and removal of CUDA Toolkits < 10.2 and devices with CC < 5.3
Deprecation and removal of CUDA Toolkits < 11.0 and devices with CC < 5.3
=========================================================================

- Support for CUDA toolkits less than 10.2 was deprecated and removed.
- Support for CUDA toolkits less than 11.0 has been removed.
- Support for devices with Compute Capability < 5.3 is deprecated and will be
removed in the future.

Expand All @@ -302,11 +302,12 @@ Recommendations

- For devices of Compute Capability 3.0 and 3.2, Numba 0.55.1 or earlier will
be required.
- CUDA toolkit 10.2 or later (ideally 11.2 or later) should be installed.
- CUDA toolkit 11.0 or later (ideally 11.2 or later) should be installed.

Schedule
--------

- In Numba 0.55.1: support for CC < 5.3 and CUDA toolkits < 10.2 was deprecated.
- In Numba 0.56: support for CC < 3.5 and CUDA toolkits < 10.2 was removed.
- In Numba 0.57: support for CC < 5.3 will be removed.
- In Numba 0.57: Support for CUDA toolkit 10.2 was removed. Support for CC < 5.3
will be removed.
2 changes: 1 addition & 1 deletion docs/source/reference/envvars.rst
Original file line number Diff line number Diff line change
Expand Up @@ -450,7 +450,7 @@ GPU support
The default compute capability (a string of the type ``major.minor``) to
target when compiling to PTX using ``cuda.compile_ptx``. The default is
5.2, which is the lowest non-deprecated compute capability in the most
recent version of the CUDA toolkit supported (10.2 at present).
recent version of the CUDA toolkit supported (11.0 at present).

.. envvar:: NUMBA_ENABLE_CUDASIM

Expand Down
9 changes: 2 additions & 7 deletions numba/cuda/cudadrv/nvvm.py
Original file line number Diff line number Diff line change
Expand Up @@ -341,16 +341,11 @@ def get_supported_ccs():

ctk_ver = f"{cudart_version[0]}.{cudart_version[1]}"
unsupported_ver = f"CUDA Toolkit {ctk_ver} is unsupported by Numba - " \
+ "10.2 is the minimum required version."
+ "11.0 is the minimum required version."

# List of supported compute capability in sorted order
if cudart_version == (0, 0):
_supported_cc = ()
elif cudart_version == (10, 2):
_supported_cc = ((3, 5), (3, 7),
(5, 0), (5, 2), (5, 3),
(6, 0), (6, 1), (6, 2),
(7, 0), (7, 2), (7, 5))
elif cudart_version == (11, 0):
_supported_cc = ((3, 5), (3, 7),
(5, 0), (5, 2), (5, 3),
Expand Down Expand Up @@ -421,7 +416,7 @@ def get_arch_option(major, minor):


MISSING_LIBDEVICE_FILE_MSG = '''Missing libdevice file.
Please ensure you have package cudatoolkit >= 10.2
Please ensure you have package cudatoolkit >= 11.0
Install package by:

conda install cudatoolkit
Expand Down
5 changes: 2 additions & 3 deletions numba/cuda/cudadrv/runtime.py
Original file line number Diff line number Diff line change
Expand Up @@ -128,9 +128,8 @@ def supported_versions(self):
if sys.platform not in ('linux', 'win32') or config.MACHINE_BITS != 64:
# Only 64-bit Linux and Windows are supported
return ()
return ((9, 2),
(10, 0), (10, 1), (10, 2),
(11, 0), (11, 1), (11, 2))
return ((11, 0), (11, 1), (11, 2), (11, 3), (11, 4), (11, 5), (11, 6),
(11, 7))


runtime = Runtime()
Expand Down
12 changes: 1 addition & 11 deletions numba/cuda/cudaimpl.py
Original file line number Diff line number Diff line change
Expand Up @@ -502,18 +502,8 @@ def ptx_fp16_hneg(context, builder, sig, args):

@lower(stubs.fp16.habs, types.float16)
def ptx_fp16_habs(context, builder, sig, args):
if cuda.runtime.get_version() < (10, 2):
# CUDA < 10.2 does not support abs.f16. For these versions, we mask
# off the sign bit to compute abs instead. We determine whether or
# not to do this based on the runtime version so that our behaviour
# is consistent with the version of NVVM we're using to go from
# NVVM IR -> PTX.
inst = 'and.b16 $0, $1, 0x7FFF;'
else:
inst = 'abs.f16 $0, $1;'

fnty = ir.FunctionType(ir.IntType(16), [ir.IntType(16)])
asm = ir.InlineAsm(fnty, inst, '=h,h')
asm = ir.InlineAsm(fnty, 'abs.f16 $0, $1;', '=h,h')
return builder.call(asm, args)


Expand Down
7 changes: 3 additions & 4 deletions numba/cuda/tests/cudadrv/test_runtime.py
Original file line number Diff line number Diff line change
Expand Up @@ -21,9 +21,8 @@ def set_visible_devices_and_check(q):
if config.ENABLE_CUDASIM:
SUPPORTED_VERSIONS = (-1, -1),
else:
SUPPORTED_VERSIONS = ((9, 2),
(10, 0), (10, 1), (10, 2),
(11, 0), (11, 1), (11, 2))
SUPPORTED_VERSIONS = ((11, 0), (11, 1), (11, 2), (11, 3), (11, 4), (11, 5),
(11, 6), (11, 7))


class TestRuntime(unittest.TestCase):
Expand All @@ -36,7 +35,7 @@ def test_is_supported_version_true(self):
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)):
for v in ((10, 2), (11, 8), (12, 0)):
with patch.object(runtime, 'get_version', return_value=v):
self.assertFalse(runtime.is_supported_version())

Expand Down
5 changes: 1 addition & 4 deletions numba/cuda/tests/cudapy/test_intrinsics.py
Original file line number Diff line number Diff line change
Expand Up @@ -557,10 +557,7 @@ def test_habs_scalar(self):
def test_habs_ptx(self):
args = (f2[:], f2)
ptx, _ = compile_ptx(simple_habs_scalar, args, cc=(5, 3))
if cuda.runtime.get_version() < (10, 2):
self.assertRegex(ptx, r'and\.b16.*0x7FFF;')
else:
self.assertIn('abs.f16', ptx)
self.assertIn('abs.f16', ptx)

@skip_unless_cc_53
def test_fp16_comparison(self):
Expand Down