Skip to content

CUDA: Add CUDA 11.8 / Hopper support and required fixes #8579

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
esc merged 5 commits into from
Dec 3, 2022
Merged

CUDA: Add CUDA 11.8 / Hopper support and required fixes #8579

Show file tree
Hide file tree
Changes from 2 commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
4a9672f
CUDA: Fix supported CC logic
gmarkall Nov 4, 2022
f7adcec
CUDA tests: always add the NVVM IR version
gmarkall Nov 4, 2022
e3e7d82
NVVM: Don't cache supported CCs at global scope
gmarkall Nov 14, 2022
c4e4081
Clarify that supported CC bounds are inclusive
gmarkall Nov 14, 2022
3e7ebb4
NVVM: Don't use hardcoded version limits
gmarkall Nov 14, 2022
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
78 changes: 48 additions & 30 deletions numba/cuda/cudadrv/nvvm.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -343,51 +343,69 @@ def get_log(self):
return ''


COMPUTE_CAPABILITIES = (
(3, 5), (3, 7),
(5, 0), (5, 2), (5, 3),
(6, 0), (6, 1), (6, 2),
(7, 0), (7, 2), (7, 5),
(8, 0), (8, 6), (8, 7), (8, 9),
(9, 0)
)

# Maps CTK version -> (min supported cc, max supported cc)

@stuartarchibald stuartarchibald Nov 7, 2022

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Are the bounds inclusive/exclusive, should this be clarified?

@gmarkall gmarkall Nov 14, 2022

Copy link
Copy Markdown
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Now clarified.

CTK_SUPPORTED = {
(11, 0): ((3, 5), (8, 0)),
(11, 1): ((3, 5), (8, 6)),
(11, 2): ((3, 5), (8, 6)),
(11, 3): ((3, 5), (8, 6)),
(11, 4): ((3, 5), (8, 7)),
(11, 5): ((3, 5), (8, 7)),
(11, 6): ((3, 5), (8, 7)),
(11, 7): ((3, 5), (8, 7)),
(11, 8): ((3, 5), (9, 0)),
}


def ccs_supported_by_ctk(ctk_version):
try:
# For supported versions, we look up the range of supported CCs
min_cc, max_cc = CTK_SUPPORTED[ctk_version]
return tuple([cc for cc in COMPUTE_CAPABILITIES
if min_cc <= cc <= max_cc])
except KeyError:
# For unsupported CUDA toolkit versions, all we can do is assume all
# non-deprecated versions we are aware of are supported.
return tuple([cc for cc in COMPUTE_CAPABILITIES if cc >= (5, 3)])


_supported_cc = None


def get_supported_ccs():
# Attempt to return cached list

@stuartarchibald stuartarchibald Nov 7, 2022

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

If there's an in-process cache need, should something like functools.cache be used opposed to module globals?

@gmarkall gmarkall Nov 7, 2022

Copy link
Copy Markdown
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

That would be way better than the current bodge!

@gmarkall gmarkall Nov 14, 2022

Copy link
Copy Markdown
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Actually it doesn't even need to be cached, this just makes no sense. The supported CCs is a property of the NVVM version, so I'm going to just make supported_ccs a property of the NVVM instance, because that just makes more sense anyway (and implicitly "caches" the values too because it's a singleton).

I think there could be an argument for the logic in get_supported_ccs() and its callees to be part of the NVVM class, but I think that functionality will also be common with NVPTX so I don't want to hoist the code into NVVM for now.

global _supported_cc

if _supported_cc:
if _supported_cc is not None:
return _supported_cc

try:
from numba.cuda.cudadrv.runtime import runtime
cudart_version = runtime.get_version()
except: # noqa: E722
# The CUDA Runtime may not be present
cudart_version = (0, 0)

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

# List of supported compute capability in sorted order
if cudart_version == (0, 0):
# We can't support anything if there's an error getting the runtime
# version (e.g. if it's not present or there's another issue)
_supported_cc = ()
elif cudart_version == (11, 0):
_supported_cc = ((3, 5), (3, 7),
(5, 0), (5, 2), (5, 3),
(6, 0), (6, 1), (6, 2),
(7, 0), (7, 2), (7, 5),
(8, 0))
elif cudart_version > (11, 0):
_supported_cc = ((3, 5), (3, 7),
(5, 0), (5, 2), (5, 3),
(6, 0), (6, 1), (6, 2),
(7, 0), (7, 2), (7, 5),
(8, 0), (8, 6))
elif cudart_version > (11, 4):
_supported_cc = ((3, 5), (3, 7),
(5, 0), (5, 2), (5, 3),
(6, 0), (6, 1), (6, 2),
(7, 0), (7, 2), (7, 5),
(8, 0), (8, 6), (8, 7))
else:
return _supported_cc

# Ensure the minimum CTK version requirement is met
if cudart_version < (11, 0):
_supported_cc = ()
ctk_ver = f"{cudart_version[0]}.{cudart_version[1]}"
unsupported_ver = (f"CUDA Toolkit {ctk_ver} is unsupported by Numba - "
"11.0 is the minimum required version.")

@stuartarchibald stuartarchibald Nov 7, 2022

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Should the minimum version be specified as a constant somewhere (maybe a module global, or perhaps even all this encoding of supported CCs etc ought to go into it's own module or YAML file?) and then code be updated in terms of that constant?

@gmarkall gmarkall Nov 14, 2022

Copy link
Copy Markdown
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think instead of hardcoding the minimum CTK required, it can be derived from the dict of supported CCs, then it doesn't need to be stored separately anywhere. Additionally, I'm going to use NUMBA_CUDA_DEFAULT_PTX_CC for the default minimum for unrecognised toolkits instead of hardcoding (5, 3) above. This will be more robust in that a constant doesn't need updating in two places, and if there's some future toolkit that doesn't even support 5.3, the user can workaround it by setting the environment variable NUMBA_CUDA_DEFAULT_PTX_CC=6.1 (or whatever) to match the minimum for their toolkit.

The CC support matrix for NVPTX will differ, so I think keeping this encoding of supported CCs in nvvm.py is logical, and keeps complexity to a minimum instead of adding additional logic / dependencies.

warnings.warn(unsupported_ver)
return _supported_cc

_supported_cc = ccs_supported_by_ctk(cudart_version)
return _supported_cc


Expand Down
1 change: 1 addition & 0 deletions numba/cuda/tests/cudadrv/test_inline_ptx.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,6 +10,7 @@ class TestCudaInlineAsm(ContextResettingTestCase):
def test_inline_rsqrt(self):
mod = ir.Module(__name__)
mod.triple = 'nvptx64-nvidia-cuda'
nvvm.add_ir_version(mod)
fnty = ir.FunctionType(ir.VoidType(), [ir.PointerType(ir.FloatType())])
fn = ir.Function(mod, fnty, 'cu_rsqrt')
bldr = ir.IRBuilder(fn.append_basic_block('entry'))
Expand Down
41 changes: 22 additions & 19 deletions numba/cuda/tests/cudadrv/test_nvvm_driver.py
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,8 +1,7 @@
import warnings

from llvmlite import ir
from numba.cuda.cudadrv.nvvm import (llvm_to_ptx, set_cuda_kernel,
get_arch_option, get_supported_ccs)
from numba.cuda.cudadrv import nvvm
from ctypes import c_size_t, c_uint64, sizeof
from numba.cuda.testing import unittest
from numba.cuda.cudadrv.nvvm import LibDevice, NvvmError, NVVM
Expand All @@ -15,7 +14,8 @@
class TestNvvmDriver(unittest.TestCase):
def get_nvvmir(self):
if NVVM().is_nvvm70:
metadata = metadata_nvvm70
versions = NVVM().get_ir_version()
metadata = metadata_nvvm70 % versions
Comment on lines +17 to +18

@stuartarchibald stuartarchibald Nov 7, 2022

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Perhaps update to use .format() or an f-string?

@gmarkall gmarkall Nov 7, 2022

Copy link
Copy Markdown
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I decided against that because it makes the IR less legible due to the presence of { and }.

else:
metadata = metadata_nvvm34

Expand All @@ -25,21 +25,22 @@ def get_nvvmir(self):

def test_nvvm_compile_simple(self):
nvvmir = self.get_nvvmir()
ptx = llvm_to_ptx(nvvmir).decode('utf8')
ptx = nvvm.llvm_to_ptx(nvvmir).decode('utf8')
self.assertTrue('simple' in ptx)
self.assertTrue('ave' in ptx)

def test_nvvm_from_llvm(self):
m = ir.Module("test_nvvm_from_llvm")
m.triple = 'nvptx64-nvidia-cuda'
nvvm.add_ir_version(m)
Comment on lines 33 to +35

@stuartarchibald stuartarchibald Nov 7, 2022

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Is this 3 line "setup" for a CUDA-target-LLVM-module sufficiently common to make it worth creating a utility function? The main concern here is that the m.triple specification is likely always necessary whereas the nvvm.add_ir_version(m) appears not to be? I suppose the underlying question is "what is the need driving the addition of the add_ir_version()" call?

@gmarkall gmarkall Nov 14, 2022

Copy link
Copy Markdown
Member Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This does happen for new modules created by the JITCUDACodegen, but I didn't use that here because it's part of the compiler / "cudapy", rather than the driver, which is under test in these cases.

fty = ir.FunctionType(ir.VoidType(), [ir.IntType(32)])
kernel = ir.Function(m, fty, name='mycudakernel')
bldr = ir.IRBuilder(kernel.append_basic_block('entry'))
bldr.ret_void()
set_cuda_kernel(kernel)
nvvm.set_cuda_kernel(kernel)

m.data_layout = NVVM().data_layout
ptx = llvm_to_ptx(str(m)).decode('utf8')
ptx = nvvm.llvm_to_ptx(str(m)).decode('utf8')
self.assertTrue('mycudakernel' in ptx)
if is64bit:
self.assertTrue('.address_size 64' in ptx)
Expand All @@ -50,40 +51,42 @@ def test_nvvm_ir_verify_fail(self):
m = ir.Module("test_bad_ir")
m.triple = "unknown-unknown-unknown"
m.data_layout = NVVM().data_layout
nvvm.add_ir_version(m)
with self.assertRaisesRegex(NvvmError, 'Invalid target triple'):
llvm_to_ptx(str(m))
nvvm.llvm_to_ptx(str(m))

def _test_nvvm_support(self, arch):
compute_xx = 'compute_{0}{1}'.format(*arch)
nvvmir = self.get_nvvmir()
ptx = llvm_to_ptx(nvvmir, arch=compute_xx, ftz=1, prec_sqrt=0,
prec_div=0).decode('utf8')
ptx = nvvm.llvm_to_ptx(nvvmir, arch=compute_xx, ftz=1, prec_sqrt=0,
prec_div=0).decode('utf8')
self.assertIn(".target sm_{0}{1}".format(*arch), ptx)
self.assertIn('simple', ptx)
self.assertIn('ave', ptx)

def test_nvvm_support(self):
"""Test supported CC by NVVM
"""
for arch in get_supported_ccs():
for arch in nvvm.get_supported_ccs():
self._test_nvvm_support(arch=arch)

def test_nvvm_warning(self):
m = ir.Module("test_nvvm_warning")
m.triple = 'nvptx64-nvidia-cuda'
m.data_layout = NVVM().data_layout
nvvm.add_ir_version(m)

fty = ir.FunctionType(ir.VoidType(), [])
kernel = ir.Function(m, fty, name='inlinekernel')
builder = ir.IRBuilder(kernel.append_basic_block('entry'))
builder.ret_void()
set_cuda_kernel(kernel)
nvvm.set_cuda_kernel(kernel)

# Add the noinline attribute to trigger NVVM to generate a warning
kernel.attributes.add('noinline')

with warnings.catch_warnings(record=True) as w:
llvm_to_ptx(str(m))
nvvm.llvm_to_ptx(str(m))

self.assertEqual(len(w), 1)
self.assertIn('overriding noinline attribute', str(w[0]))
Expand All @@ -107,14 +110,14 @@ def test_nvvm_future_support(self):
class TestArchOption(unittest.TestCase):
def test_get_arch_option(self):
# Test returning the nearest lowest arch.
self.assertEqual(get_arch_option(5, 3), 'compute_53')
self.assertEqual(get_arch_option(7, 5), 'compute_75')
self.assertEqual(get_arch_option(7, 7), 'compute_75')
self.assertEqual(nvvm.get_arch_option(5, 3), 'compute_53')
self.assertEqual(nvvm.get_arch_option(7, 5), 'compute_75')
self.assertEqual(nvvm.get_arch_option(7, 7), 'compute_75')
# Test known arch.
supported_cc = get_supported_ccs()
supported_cc = nvvm.get_supported_ccs()
for arch in supported_cc:
self.assertEqual(get_arch_option(*arch), 'compute_%d%d' % arch)
self.assertEqual(get_arch_option(1000, 0),
self.assertEqual(nvvm.get_arch_option(*arch), 'compute_%d%d' % arch)
self.assertEqual(nvvm.get_arch_option(1000, 0),
'compute_%d%d' % supported_cc[-1])


Expand Down Expand Up @@ -163,7 +166,7 @@ def test_libdevice_load(self):

metadata_nvvm70 = '''
!nvvmir.version = !{!1}
!1 = !{i32 1, i32 6, i32 3, i32 0}
!1 = !{i32 %s, i32 %s, i32 %s, i32 %s}

!nvvm.annotations = !{!2}
!2 = !{void (i32*)* @simple, !"kernel", i32 1}
Expand Down
1 change: 1 addition & 0 deletions numba/cuda/tests/nocuda/test_nvvm.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -62,6 +62,7 @@ def test_nvvm_accepts_encoding(self):
bytearray(range(256)))
m = ir.Module()
m.triple = 'nvptx64-nvidia-cuda'
nvvm.add_ir_version(m)
gv = ir.GlobalVariable(m, c.type, "myconstant")
gv.global_constant = True
gv.initializer = c
Expand Down