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_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 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 _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 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 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(): 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 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')