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)
def test_create_program(self): brig_file = get_brig_file() symbol = '&__vector_copy_kernel' brig_module = BrigModule.from_file(brig_file) program = Program() program.add_module(brig_module) code = program.finalize(self.gpu.isa) ex = Executable() ex.load(self.gpu, code) ex.freeze() sym = ex.get_symbol(self.gpu, symbol) self.assertGreater(sym.kernarg_segment_size, 0)
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_from_file(self): brig_file = get_brig_file() brig_module = BrigModule.from_file(brig_file) self.assertGreater(len(brig_module), 0)