def test_nvvm_memset_fixup_missing_align(self): """ We require alignment to be specified as a parameter attribute to the dest argument of a memset. """ with self.assertRaises(ValueError) as e: nvvm.llvm39_to_34_ir(missing_align) self.assertIn(str(e.exception), "No alignment attribute found on memset dest")
def test_nvvm_memset_fixup(self): """ Test llvm.memset changes in llvm7. In LLVM7 the alignment parameter can be implicitly provided as an attribute to pointer in the first argument. """ def foo(x): # Triggers a generation of llvm.memset for i in range(x.size): x[i] = 0 cukern = compile_kernel(foo, args=(types.int32[::1], ), link=()) original = cukern._func.ptx.llvmir self.assertIn("call void @llvm.memset", original) fixed = nvvm.llvm39_to_34_ir(original) self.assertIn("call void @llvm.memset", fixed) # Check original IR for ln in original.splitlines(): if 'call void @llvm.memset' in ln: # Missing i32 4 in the 2nd last argument self.assertRegexpMatches( ln, r'i64 %\d+, i1 false\)'.replace(' ', r'\s+'), ) # Check fixed IR for ln in fixed.splitlines(): if 'call void @llvm.memset' in ln: # The i32 4 is the alignment self.assertRegexpMatches( ln, r'i32 4, i1 false\)'.replace(' ', r'\s+'), )
def test_nvvm_memset_fixup(self): """ Test llvm.memset changes in llvm7. In LLVM7 the alignment parameter can be implicitly provided as an attribute to pointer in the first argument. """ def foo(x): # Triggers a generation of llvm.memset for i in range(x.size): x[i] = 0 cukern = compile_kernel(foo, args=(types.int32[::1],), link=()) original = cukern._func.ptx.llvmir self.assertIn("call void @llvm.memset", original) fixed = nvvm.llvm39_to_34_ir(original) self.assertIn("call void @llvm.memset", fixed) # Check original IR for ln in original.splitlines(): if 'call void @llvm.memset' in ln: # Missing i32 4 in the 2nd last argument self.assertRegexpMatches( ln, r'i64 %\d+, i1 false\)'.replace(' ', r'\s+'), ) # Check fixed IR for ln in fixed.splitlines(): if 'call void @llvm.memset' in ln: # The i32 4 is the alignment self.assertRegexpMatches( ln, r'i32 4, i1 false\)'.replace(' ', r'\s+'), )
def test_nvvm_memset_fixup(self): """ Test llvm.memset changes in llvm7. In LLVM7 the alignment parameter can be implicitly provided as an attribute to pointer in the first argument. """ fixed = nvvm.llvm39_to_34_ir(original) self.assertIn("call void @llvm.memset", fixed) for ln in fixed.splitlines(): if 'call void @llvm.memset' in ln: # The i32 4 is the alignment self.assertRegexpMatches( ln, r'i32 4, i1 false\)'.replace(' ', r'\s+'), )
def patch(self, ir): # Import here to avoid error in CUDASIM from numba.cuda.cudadrv.nvvm import llvm39_to_34_ir return llvm39_to_34_ir(ir)
def patch(self, ir): # Import here to avoid error in CUDASIM from numba.cuda.cudadrv.nvvm import llvm39_to_34_ir return llvm39_to_34_ir(ir)