def run_test():
if not test():
return False
nvvm = NVVM()
print("NVVM version", nvvm.get_version())
return nvvm.get_version() is not None
def get_ptx(self):
nvvm = NVVM()
print(nvvm.get_version())
if is64bit:
return gpu64
else:
return gpu32
def get_ptx(self):
nvvm = NVVM()
print(nvvm.get_version())
if is64bit:
return gpu64
else:
return gpu32
def run_test():
if not test():
return False
nvvm = NVVM()
print("NVVM version", nvvm.get_version())
# check pkg version matches lib pulled in
gotlib = get_cudalib('cublas')
lookfor = '9.2'
if sys.platform.startswith('win'):
# windows libs have no dot
lookfor = lookfor.replace('.', '')
return lookfor in gotlib
def run_test():
if not test():
return False
nvvm = NVVM()
print("NVVM version", nvvm.get_version())
# check pkg version matches lib pulled in
gotlib = get_cudalib('cublas')
lookfor = os.environ['PKG_VERSION']
if sys.platform.startswith('win'):
# windows libs have no dot
lookfor = lookfor.replace('.', '')
return lookfor in gotlib
def run_pass(self, state):
# Early return if NVVM 7
from numba.cuda.cudadrv.nvvm import NVVM
if NVVM().is_nvvm70:
return False
# NVVM < 7, need to check for charseq
typmap = state.typemap
def check_dtype(dtype):
if isinstance(dtype, (types.UnicodeCharSeq, types.CharSeq)):
msg = (f"{k} is a char sequence type. This type is not "
"supported with CUDA toolkit versions < 11.2. To "
"use this type, you need to update your CUDA "
"toolkit - try 'conda install cudatoolkit=11' if "
"you are using conda to manage your environment.")
raise TypingError(msg)
elif isinstance(dtype, types.Record):
for subdtype in dtype.fields.items():
# subdtype is a (name, _RecordField) pair
check_dtype(subdtype[1].type)
for k, v in typmap.items():
if isinstance(v, types.Array):
check_dtype(v.dtype)
return False
def get_ptx(self):
nvvm = NVVM()
if is64bit:
return gpu64
else:
return gpu32
def test_wrapper_has_debuginfo(self):
sig = (types.int32[::1],)
@cuda.jit(sig, debug=True, opt=0)
def f(x):
x[0] = 1
llvm_ir = f.inspect_llvm(sig)
if NVVM().is_nvvm70:
# NNVM 7.0 IR attaches a debug metadata reference to the
# definition
defines = [line for line in llvm_ir.splitlines()
if 'define void @"_ZN6cudapy' in line]
# Make sure we only found one definition
self.assertEqual(len(defines), 1)
wrapper_define = defines[0]
self.assertIn('noinline!dbg', wrapper_define)
else:
# NVVM 3.4 subprogram debuginfo refers to the definition.
# '786478' is a constant referring to a subprogram.
disubprograms = [line for line in llvm_ir.splitlines()
if '786478' in line]
# Make sure we only found one subprogram
self.assertEqual(len(disubprograms), 1)
wrapper_disubprogram = disubprograms[0]
# Check that the subprogram points to a wrapper (these are all in
# the "cudapy::" namespace).
self.assertIn('_ZN6cudapy', wrapper_disubprogram)
def get_nvvmir(self):
if NVVM().is_nvvm70:
metadata = metadata_nvvm70
else:
metadata = metadata_nvvm34
return nvvmir_generic + metadata
def _check(self, fn, sig, expect):
if not NVVM().is_nvvm70:
self.skipTest("debuginfo not generated for NVVM 3.4")
asm = self._getasm(fn, sig=sig)
re_section_dbginfo = re.compile(r"\.section\s+\.debug_info\s+{")
match = re_section_dbginfo.search(asm)
assertfn = self.assertIsNotNone if expect else self.assertIsNone
assertfn(match, msg=asm)
def check_warnings(self, warnings, warn_count):
if NVVM().is_nvvm70:
# We should not warn on NVVM 7.0.
self.assertEqual(len(warnings), 0)
else:
self.assertEqual(len(warnings), warn_count)
# Each warning should warn about not generating debug info.
for warning in warnings:
self.assertIs(warning.category, NumbaInvalidConfigWarning)
self.assertIn('debuginfo is not generated',
str(warning.message))
def check_debug_info(self, ptx):
if not NVVM().is_nvvm70:
self.skipTest('debuginfo not generated for NVVM 3.4')
# A debug_info section should exist in the PTX. Whitespace varies
# between CUDA toolkit versions.
self.assertRegex(ptx, '\\.section\\s+\\.debug_info')
# A .file directive should be produced and include the name of the
# source. The path and whitespace may vary, so we accept anything
# ending in the filename of this module.
self.assertRegex(ptx, '\\.file.*test_compiler.py"')
def run_test():
# on windows only nvvm is available to numba
if sys.platform.startswith('win'):
nvvm = NVVM()
print("NVVM version", nvvm.get_version())
return nvvm.get_version() is not None
if not test():
return False
nvvm = NVVM()
print("NVVM version", nvvm.get_version())
# check pkg version matches lib pulled in
gotlib = get_cudalib('cublas')
lookfor = os.environ['PKG_VERSION']
return lookfor in gotlib
def run_test():
# on windows only nvvm is available to numba
if sys.platform.startswith('win'):
nvvm = NVVM()
print("NVVM version", nvvm.get_version())
return nvvm.get_version() is not None
if not test():
return False
nvvm = NVVM()
print("NVVM version", nvvm.get_version())
# check pkg version matches lib pulled in
gotlib = get_cudalib('cublas')
# check cufft b/c cublas has an incorrect version in 10.1 update 1
gotlib = get_cudalib('cufft')
return bool(get_cudalib('cublas')) and bool(get_cudalib('cufft'))
def run_test():
# on windows only nvvm is available to numba
if sys.platform.startswith("win"):
nvvm = NVVM()
print("NVVM version", nvvm.get_version())
return nvvm.get_version() is not None
if not test():
return False
nvvm = NVVM()
print("NVVM version", nvvm.get_version())
extra_lib_tests = (
"cublas", # check pkg version matches lib pulled in
"cufft", # check cufft b/c cublas has an incorrect version in 10.1 update 1
"cupti", # check this is getting included
)
found_paths = [get_cudalib(lib) for lib in extra_lib_tests]
print(*zip(extra_lib_tests, found_paths), sep="\n")
return all(extra_lib_tests)
def skip_on_nvvm70(self):
if NVVM().is_nvvm70:
self.skipTest('Character sequences are permitted with NVVM 7.0')
def skip_on_nvvm34(self):
# There is no NVVM on cudasim, but we might as well run this test.
if not config.ENABLE_CUDASIM and not NVVM().is_nvvm70:
self.skipTest('Character sequences unsupported on NVVM 3.4')
