예제 #1
0
    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)
예제 #2
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, "&copy")
        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)
예제 #3
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)
예제 #4
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)
예제 #5
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, "&copy")
        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)