Example #1
0
def test_copy_texture_memory_args(drv, *args):
    drv = setup_mock(drv)

    fake_array = np.zeros(10).astype(np.float32)
    texmem_args = {'fake_tex': fake_array}

    texref = Mock()

    with pycuda.PyCudaFunctions(0) as dev:
        dev.current_module = Mock()
        dev.current_module.get_texref.return_value = texref

        dev.copy_texture_memory_args(texmem_args)

        drv.matrix_to_texref.assert_called_once_with(fake_array, texref, order="C")
        dev.current_module.get_texref.assert_called_once_with('fake_tex')

        texmem_args = {'fake_tex2': {'array': fake_array, 'filter_mode': 'linear', 'address_mode': ['border', 'clamp']}}

        dev.copy_texture_memory_args(texmem_args)
        drv.matrix_to_texref.assert_called_with(fake_array, texref, order="C")
        dev.current_module.get_texref.assert_called_with('fake_tex2')
        texref.set_filter_mode.assert_called_once_with(drv.filter_mode.LINEAR)
        texref.set_address_mode.assert_any_call(0, drv.address_mode.BORDER)
        texref.set_address_mode.assert_any_call(1, drv.address_mode.CLAMP)
Example #2
0
def test_benchmark(drv, *args):
    drv = setup_mock(drv)

    drv.Event.return_value.time_since.return_value = 0.1

    with pycuda.PyCudaFunctions(0) as dev:
        res = dev.benchmark(dummy_func, [1, 2], (1, 2), (1, 2))
        assert res["time"] > 0

        assert dev.context.synchronize.call_count > 1
        assert drv.Event.return_value.synchronize.call_count == dev.iterations
        assert drv.Event.return_value.record.call_count == 2*dev.iterations
        assert drv.Event.return_value.time_since.call_count == dev.iterations
Example #3
0
def test_copy_constant_memory_args(drv, *args):
    drv = setup_mock(drv)

    fake_array = np.zeros(10).astype(np.float32)
    cmem_args = {'fake_array': fake_array}

    with pycuda.PyCudaFunctions(0) as dev:
        dev.current_module = Mock()
        dev.current_module.get_global.return_value = ['get_global']

        dev.copy_constant_memory_args(cmem_args)

        drv.memcpy_htod.assert_called_once_with('get_global', fake_array)
        dev.current_module.get_global.assert_called_once_with('fake_array')
def test_ready_argument_list():

    size = 1000
    a = np.int32(75)
    b = np.random.randn(size).astype(np.float32)
    c = np.zeros_like(b)

    arguments = [c, a, b]

    with pycuda.PyCudaFunctions(0) as dev:
        gpu_args = dev.ready_argument_list(arguments)

        assert isinstance(gpu_args[0], pycuda.driver.DeviceAllocation)
        assert isinstance(gpu_args[1], np.int32)
        assert isinstance(gpu_args[2], pycuda.driver.DeviceAllocation)
Example #5
0
def test_ready_argument_list(drv, *args):
    drv = setup_mock(drv)

    size = 5
    a = np.int32(75)
    b = np.random.randn(size).astype(np.float32)
    arguments = [a, b]

    with pycuda.PyCudaFunctions(0) as dev:
        gpu_args = dev.ready_argument_list(arguments)

    print(drv.mock_calls)
    print(gpu_args)

    drv.mem_alloc.assert_called_once_with(20)
    drv.memcpy_htod.assert_called_once_with('mem_alloc', b)

    assert isinstance(gpu_args[0], np.int32)
def test_compile():

    kernel_string = """
    __global__ void vector_add(float *c, float *a, float *b, int n) {
        int i = blockIdx.x * blockDim.x + threadIdx.x;
        if (i<n) {
            c[i] = a[i] + b[i];
        }
    }
    """

    kernel_name = "vector_add"
    kernel_sources = KernelSource(kernel_name, kernel_string, "cuda")
    kernel_instance = KernelInstance(kernel_name, kernel_sources, kernel_string, [], None, None, dict(), [])
    with pycuda.PyCudaFunctions(0) as dev:
        try:
            dev.compile(kernel_instance)
        except Exception as e:
            pytest.fail("Did not expect any exception:" + str(e))
Example #7
0
def test_compile(drv, *args):

    # setup mocked stuff
    drv = setup_mock(drv)
    with pycuda.PyCudaFunctions(0) as dev:
        dev.source_mod = Mock()
        dev.source_mod.return_value.get_function.return_value = 'func'

        # call compile
        kernel_string = "__global__ void vector_add()"
        kernel_name = "vector_add"
        kernel_sources = KernelSource(kernel_name, kernel_string, "cuda")
        kernel_instance = KernelInstance(kernel_name, kernel_sources, kernel_string, [], None, None, dict(), [])
        func = dev.compile(kernel_instance)

        # verify behavior
        assert dev.source_mod.call_count == 1
        assert dev.current_module is dev.source_mod.return_value
        assert func == 'func'

        assert kernel_string == list(dev.source_mod.mock_calls[0])[1][0]
        optional_args = list(dev.source_mod.mock_calls[0])[2]
        assert optional_args['code'] == 'sm_55'
        assert optional_args['arch'] == 'compute_55'
def test_benchmark():
    with pycuda.PyCudaFunctions(0) as dev:
        args = [1, 2]
        res = dev.benchmark(dummy_func, args, (1, 2), (1, 2))
        assert res["time"] > 0
        assert len(res["times"]) == dev.iterations