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)