def test_current_queue_feature_2():
set_current_queue(None)
queue1 = get_current_queue()
# use get_devices() to get list with available_devices where index corresponds to device id
set_current_queue(queue1)
thread1_reused = get_current_queue()
assert hash(queue1) == hash(thread1_reused)
set_current_queue(create_queue(device_id=0))
thread2 = get_current_queue()
assert hash(queue1) != hash(thread2)
set_current_queue(queue1)
thread1_reused = get_current_queue()
assert hash(queue1) == hash(thread1_reused)
def eval_code(emulate=False):
data = to_device(np.array([0]).astype(Types.char))
knl = Kernel('knl_test_for_loop', {
'data': Global(data)
},
"""
${header}{
data[0]+=i;
}
""",
replacements={
'header': header
},
global_size=data.shape).compile(emulate=emulate)
knl()
get_current_queue().finish()
res = knl.data.get()
return res
def test_bit_shift(
): # todo use https://numpy.org/doc/stable/reference/generated/numpy.ndarray.ctypes.html
data = np.array([0, 0, 0, 0]).astype(Types.char)
knl = Kernel('knl_bit_packing', {'data': data},
"""
uchar a = 5;
uchar b = 3;
uchar c = (a << 4) | b;
data[0] = (c & 0xF0) >> 4;
data[1] = c & (0x0F);
""",
global_size=data.shape)
prog = Program(kernels=[knl])
knl_cl = prog.compile().knl_bit_packing
knl_py = prog.compile(emulate=True).knl_bit_packing
knl_cl()
get_current_queue().finish()
res_cl = knl_cl.data.get()
knl_py()
res_py = knl_cl.data.get()
assert np.all(res_cl == res_py)
def test_vector_types(
): # todo use https://numpy.org/doc/stable/reference/generated/numpy.ndarray.ctypes.html
data = np.zeros((10, )).astype(Types.char2)
knl = Kernel('knl_vector_types', {'data': data},
"""
char2 a = (char2)(4,2);
char2 b = (char2)(1,2);
data[0] = a;
data[1] = b;
data[2] = a + b;
data[3] = a * b;
data[4] = a - b;
data[5] = a / b;
""",
global_size=data.shape)
knl_cl = knl.compile()
knl_py = knl.compile(emulate=True)
knl_cl()
get_current_queue().finish()
res_cl = knl_cl.data.get()
knl_py()
res_py = knl_py.data.get()
assert np.all(res_cl == res_py)
def test_debug_kernel_with_complex_numbers():
buff_cl = zeros((10, 1), Types.cfloat)
queue = get_current_queue()
# compute result with opencl
program = pyopencl_program(queue.context,
str(rendered_template_cplx_float)).build()
some_operation = program.all_kernels()[0]
# some_operation.set_scalar_arg_dtypes([None, ClTypes.cfloat])
some_operation(queue, buff_cl.shape, None, buff_cl.data,
np.dtype(Types.cfloat).type(3 + 1j))
buff_py = zeros((10, 1), Types.cfloat)
code_py = unparse_c_code_to_python(rendered_template_cplx_float)
module = create_py_file_and_load_module(
code_py, str(path_py_cl.joinpath('debug_component_complex_numbers')))
module.some_operation(buff_py.shape, None, buff_py,
np.dtype(Types.cfloat).type(3 + 1j))
assert np.all(buff_py.get() == buff_cl.get())
def test_debug_kernel():
buff_cl = zeros((10, 1), Types.short)
# compute result with opencl
queue = get_current_queue()
program = pyopencl_program(queue.context, str(rendered_template)).build()
some_operation = program.all_kernels()[0]
# some_operation(init.queue, buff_cl.shape, None,
# buff_cl.data, np.dtype(ClTypes.short).type(3))
# alternative:
some_operation.set_scalar_arg_dtypes([None, Types.short])
some_operation(queue, buff_cl.shape, None, buff_cl.data, 3)
buff_py = zeros((10, 1), Types.short)
code_py = unparse_c_code_to_python(rendered_template)
module = create_py_file_and_load_module(
code_py, str(path_py_cl.joinpath('debug_component')))
module.some_operation(buff_py.shape, None, buff_py,
np.dtype(Types.short).type(3))
assert np.all(buff_py.get() == 11)
assert np.all(buff_py.get() == buff_cl.get())
},
"""
__local int shared[2];
ary[get_global_id(0)] = parent(ary, shared);
""",
global_size=ary.shape)
prog = Program([func_nested, func_parent], [knl])
prog_py = prog.compile(emulate=True)
prog_cl = prog.compile(emulate=False)
prog_py.some_knl()
ary_py = ary.get()
ary.set(ary_np)
prog_cl.some_knl()
ary_cl = ary.get()
get_current_queue().finish()
assert np.allclose(ary_py, np.array([2, 1]))
assert np.allclose(ary_py, ary_cl)
def test_macro_with_arguments():
defines = {
'FUNC(a,b,c)': '{ int tmp = c(a-b); a += b + tmp; }'
} # this is a macro with arguments
ary = zeros((2, ), Types.int)
func_add_two = Function('add_two', {'a': Scalar(Types.int)},
'return a + 2;',
returns=Types.int)
knl = Kernel('knl_macro_func', {'ary': Global(ary)},
"""
int a = 1;
def test_current_queue_feature_1():
# whenever create_thread() is called this is set as the current thread. Be careful, that Thread() leads to a new
# context/queue.
queue = create_queue()
assert hash(get_current_queue()) == hash(queue)