def _test_template(self, nbtype, src):
dtype = np.dtype(str(nbtype))
dst = np.zeros(1, dtype=dtype)
src = dtype.type(src)
arytype = nbtype[::1]
kernel = compiler.compile_kernel(assign_value, [arytype, nbtype])
kernel[1, 1](dst, src)
self.assertEqual(dst[0], src)
def test_loading_from_memory(self):
arytype = types.float32[:]
kernel = compiler.compile_kernel(copy_kernel_1d, [arytype] * 2)
# Load BRIG memory
symbol = '&{0}'.format(kernel.entry_name)
brig_module = BrigModule(kernel.binary)
self._check(brig_module, symbol)
def test_hsa_kernel(self):
src = np.arange(1024, dtype=np.float32)
dst = np.zeros_like(src)
# Compiler kernel
arytype = types.float32[::1]
kernel = compiler.compile_kernel(copy_kernel_1d, [arytype] * 2)
# Run kernel
kernel[src.size // 256, 256](dst, src)
np.testing.assert_equal(src, dst)
def test_loading_from_file(self):
arytype = types.float32[:]
kernel = compiler.compile_kernel(copy_kernel_1d, [arytype] * 2)
# Write the brig file out
brig_file = tempfile.NamedTemporaryFile(delete=False)
with brig_file:
brig_file.write(kernel.binary)
# Load BRIG file
symbol = '&{0}'.format(kernel.entry_name)
brig_module = BrigModule.from_file(brig_file.name)
# Cleanup
os.unlink(brig_file.name)
self._check(brig_module, symbol)
def test_device_function(self):
src = np.arange(10, dtype=np.int32)
dst = np.zeros_like(src)
arytype = types.int32[::1]
devfn = compiler.compile_device(udt_devfunc, arytype.dtype,
[arytype, types.intp])
def udt_devfunc_caller(dst, src):
i = hsa.get_global_id(0)
if i < dst.size:
dst[i] = devfn(src, i)
kernel = compiler.compile_kernel(udt_devfunc_caller,
[arytype, arytype])
kernel[src.size, 1](dst, src)
np.testing.assert_equal(dst, src)
def test_copy_kernel_1d(self):
arytype = types.float32[:]
kernel = compiler.compile_kernel(copy_kernel_1d, [arytype] * 2)
self.assertIn("prog kernel &{0}".format(kernel.entry_name),
kernel.assembly)