def _check(self, brig_module, symbol):
prog = Program()
prog.add_module(brig_module)
code = prog.finalize(self.gpu.isa)
ex = Executable()
ex.load(self.gpu, code)
ex.freeze()
sym = ex.get_symbol(self.gpu, symbol)
self.assertTrue(sym.kernel_object)
self.assertGreater(sym.kernarg_segment_size, 0)
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_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_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)
