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)
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
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)
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))
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