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)
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)
def test_register(self): src = np.random.random(1024).astype(np.float32) hsa.hsa_memory_register(src.ctypes.data, src.nbytes) hsa.hsa_memory_deregister(src.ctypes.data, src.nbytes)