Skip to content

Instantly share code, notes, and snippets.

@arghdos arghdos/masktest.py Secret

Created Mar 19, 2018
Embed
What would you like to do?
masktest
import loopy as lp
import pyopencl as cl
import numpy as np
from loopy.kernel.data import temp_var_scope as scopes
ctx = cl.create_some_context()
queue = cl.CommandQueue(ctx)
# first broken case -- incorrect promotion of temporaries to vector dtypes
knl = lp.make_kernel(
'{[i,j]: 0 <= i,j < 12}',
"""
for j
for i
<> test = mask[i]
if test
a[i, j] = 1
end
end
end
""",
[lp.GlobalArg('a', shape=(12, 12)),
lp.TemporaryVariable('mask', shape=(12,), initializer=np.array(
np.arange(12) >= 6, dtype=np.int), read_only=True,
scope=scopes.GLOBAL)])
knl = lp.split_iname(knl, 'j', 4, inner_tag='vec')
knl = lp.split_array_axis(knl, 'a', 1, 4)
knl = lp.tag_array_axes(knl, 'a', 'N1,N0,vec')
print(lp.generate_code_v2(knl).device_code())
ans = np.zeros((12, 3, 4))
ans[6:, :, :] = 1
assert np.array_equal(knl(queue, a=np.zeros((12, 3, 4)))[1][0], ans)
$ python masktest.py
Choose platform:
[0] <pyopencl.Platform 'Intel(R) OpenCL' at 0xbfcbf0>
[1] <pyopencl.Platform 'Portable Computing Language' at 0x7efe61bdc0f0>
[2] <pyopencl.Platform 'AMD Accelerated Parallel Processing' at 0x7efe5d100430>
Choice [0]:
Set the environment variable PYOPENCL_CTX='' to avoid being asked again.
/home/ncurtis/.local/lib/python2.7/site-packages/loopy/target/pyopencl.py:146: LoopyAdvisory: in kernel loopy_kernel: No device parameter was passed to the PyOpenCLTarget. Perhaps you want to pass a device to benefit from additional checking. (add 'no_device_in_pre_codegen_checks' to silenced_warnings kernel attribute to disable)
"additional checking.", LoopyAdvisory)
/home/ncurtis/.local/lib/python2.7/site-packages/loopy/target/pyopencl.py:452: UserWarning: loopy_kernel: device not supplied to PyOpenCLTarget--workarounds for broken OpenCL implementations (such as those relating to complex numbers) may not be enabled when needed. To avoid this, pass target=lp.PyOpenCLTarget(dev) when creating the kernel.
.format(knl_name=kernel.name))
#define lid(N) ((int) get_local_id(N))
#define gid(N) ((int) get_group_id(N))
__constant long const mask[12] = { 0, 0, 0, 0, 0, 0, 1, 1, 1, 1, 1, 1 };
__kernel void __attribute__ ((reqd_work_group_size(1, 1, 1))) loopy_kernel(__global int4 *__restrict__ a)
{
long4 test;
for (int j_outer = 0; j_outer <= 2; ++j_outer)
for (int i = 0; i <= 11; ++i)
{
test = mask[i];
if (test)
a[3 * i + j_outer] = 1;
}
}
1 error generated.
Traceback (most recent call last):
File "masktest.py", line 36, in <module>
assert np.array_equal(knl(queue, a=np.zeros((12, 3, 4)))[1][0], ans)
File "/home/ncurtis/.local/lib/python2.7/site-packages/loopy/kernel/__init__.py", line 1260, in __call__
return kex(*args, **kwargs)
File "/home/ncurtis/.local/lib/python2.7/site-packages/loopy/target/pyopencl_execution.py", line 347, in __call__
kernel_info = self.kernel_info(self.arg_to_dtype_set(kwargs))
File "/home/ncurtis/.local/lib/python2.7/site-packages/pytools/__init__.py", line 578, in wrapper
result = function(obj, *args, **kwargs)
File "/home/ncurtis/.local/lib/python2.7/site-packages/loopy/target/pyopencl_execution.py", line 302, in kernel_info
.build(options=kernel.options.cl_build_options))
File "/home/ncurtis/.local/lib/python2.7/site-packages/pyopencl/__init__.py", line 462, in build
options_bytes=options_bytes, source=self._source)
File "/home/ncurtis/.local/lib/python2.7/site-packages/pyopencl/__init__.py", line 506, in _build_and_catch_errors
raise err
pyopencl.cffi_cl.RuntimeError: clBuildProgram failed: BUILD_PROGRAM_FAILURE -
Build on <pyopencl.Device 'Intel(R) Xeon(R) CPU E5-4640 v2 @ 2.20GHz' on 'Intel(R) OpenCL' at 0xc052d8>:
Compilation started
1:17:7: error: statement requires expression of scalar type ('long4' (vector of 4 'long' values) invalid)
Compilation failed
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
You can’t perform that action at this time.