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_brig(self):
# Genreate BRIG
hlcmod = hlc.Module()
hlcmod.load_llvm(SPIR_SAMPLE)
brig = hlcmod.finalize().brig
# Check the first 8 bytes for the magic string
self.assertEqual(brig[:8].decode('latin1'), 'HSA BRIG')
# Compile
from numba.hsa.hsadrv.driver import BrigModule, Program, hsa, Executable
agent = hsa.components[0]
brigmod = BrigModule(brig)
prog = Program()
prog.add_module(brigmod)
code = prog.finalize(agent.isa)
ex = Executable()
ex.load(agent, code)
ex.freeze()
sym = ex.get_symbol(agent, "©")
self.assertNotEqual(sym.kernel_object, 0)
self.assertGreater(sym.kernarg_segment_size, 0)
# Execute
import ctypes
import numpy as np
sig = hsa.create_signal(1)
kernarg_region = [r for r in agent.regions if r.supports_kernargs][0]
kernarg_types = (ctypes.c_void_p * 2)
kernargs = kernarg_region.allocate(kernarg_types)
src = np.random.random(1).astype(np.float32)
dst = np.zeros_like(src)
kernargs[0] = src.ctypes.data
kernargs[1] = dst.ctypes.data
hsa.hsa_memory_register(src.ctypes.data, src.nbytes)
hsa.hsa_memory_register(dst.ctypes.data, dst.nbytes)
hsa.hsa_memory_register(ctypes.byref(kernargs),
ctypes.sizeof(kernargs))
queue = agent.create_queue_single(32)
queue.dispatch(sym, kernargs, workgroup_size=(1, ), grid_size=(1, ))
np.testing.assert_equal(dst, src)
