def test_compile(drv, *args):
# setup mocked stuff
drv = setup_mock(drv)
with cuda.CudaFunctions(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_complies_fortran_function_with_module():
kernel_string = """
module my_fancy_module
use iso_c_binding
contains
function my_test_function() result(time)
use iso_c_binding
real (c_float) :: time
time = 42.0
end function my_test_function
end module my_fancy_module
"""
kernel_name = "my_test_function"
kernel_sources = KernelSource(kernel_string, "C")
kernel_instance = KernelInstance(kernel_name, kernel_sources, kernel_string, [], None, None, dict(), [])
cfunc = CFunctions(compiler="gfortran")
func = cfunc.compile(kernel_instance)
result = cfunc.run_kernel(func, [], (), ())
assert np.isclose(result, 42.0)
def test_compile(npct, subprocess):
kernel_string = "this is a fake C program"
kernel_name = "blabla"
kernel_sources = KernelSource(kernel_string, "C")
kernel_instance = KernelInstance(kernel_name, kernel_sources, kernel_string, [], None, None, dict(), [])
cfunc = CFunctions()
f = cfunc.compile(kernel_instance)
print(subprocess.mock_calls)
print(npct.mock_calls)
print(f)
assert len(subprocess.mock_calls) == 6
assert npct.load_library.called == 1
args, _ = npct.load_library.call_args_list[0]
filename = args[0]
print('filename=' + filename)
# check if temporary files are cleaned up correctly
import os.path
assert not os.path.isfile(filename + ".cu")
assert not os.path.isfile(filename + ".o")
assert not os.path.isfile(filename + ".so")
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_sources = KernelSource(kernel_string, "cuda")
kernel_instance = KernelInstance("vector_add", kernel_sources,
kernel_string, [], None, None, dict(), [])
dev = cuda.CudaFunctions(0)
try:
dev.compile(kernel_instance)
except Exception as e:
pytest.fail("Did not expect any exception:" + str(e))
def test_compile():
original_kernel = """
__kernel void sum(__global const float *a_g, __global const float *b_g, __global float *res_g) {
int gid = get_global_id(0);
__local float test[shared_size];
test[0] = a_g[gid];
res_g[gid] = test[0] + b_g[gid];
}
"""
kernel_sources = KernelSource("sum", original_kernel, "opencl")
kernel_string = original_kernel.replace("shared_size", str(1024))
kernel_instance = KernelInstance("sum", kernel_sources, kernel_string, [], None, None, dict(), [])
with opencl.OpenCLFunctions(0) as dev:
func = dev.compile(kernel_instance)
assert isinstance(func, pyopencl.Kernel)
def test_compile_detects_device_code(npct, subprocess):
kernel_string = "this code clearly contains device code __global__ kernel(float* arg){ return; }"
kernel_name = "blabla"
kernel_sources = KernelSource(kernel_string, "C")
kernel_instance = KernelInstance(kernel_name, kernel_sources, kernel_string, [], None, None, dict(), [])
cfunc = CFunctions()
cfunc.compile(kernel_instance)
print(subprocess.check_call.call_args_list)
# assert the filename suffix used for source compilation is .cu
dot_cu_used = False
for call in subprocess.check_call.call_args_list:
args, kwargs = call
args = args[0]
print(args)
if args[0] == 'nvcc' and args[1] == '-c':
assert args[2][-3:] == '.cu'
dot_cu_used = True
assert dot_cu_used