Hi
Since the latest beignet release (1.1.1) provided by debian does not
support double precision floating point, I ended up compiling from
source and installed the latest git code locally with
EXPERIMENTAL_DOUBLE enabled.
Successfully, clinfo does detect the cl_khr_fp64 extension. However, the
pyopencl framework I'm working with expects the extension to be included
as core functionality with OpenCL 1.2+. As an example, the attached
minimal python code (test-double-record.py) attempts to map a
numpy.float64 type to an OpenCL C struct, and in doing so, generates a
minimal kernel to test the struct (tmpvRXvhp.cl). As seen, the kernel
only enables cl_khr_fp64 if using OpenCL < 1.2, but since I'm using
Beignet (opencl 1.2), and cl_khr_fp64 is not part of core functionality,
I'm met with the following error as the generated kernel fails to compile:
Traceback (most recent call last):
File "./test-double-record.py", line 14, in <module>
numpy.dtype([('double_0', numpy.float64)]))
File "<decorator-gen-3>", line 2, in match_dtype_to_c_struct
File "/usr/lib/python2.7/dist-packages/pytools/__init__.py", line
430, in _deco
result = func(*args)
File "/usr/lib/python2.7/dist-packages/pyopencl/tools.py", line
603, in match_dtype_to_c_struct
knl = prg.build(devices=[device]).get_size_and_offsets
File "/usr/lib/python2.7/dist-packages/pyopencl/__init__.py", line
213, in build
options=options, source=self._source)
File "/usr/lib/python2.7/dist-packages/pyopencl/__init__.py", line
253, in _build_and_catch_errors
raise err
pyopencl.RuntimeError: clBuildProgram failed: build program failure -
Build on <pyopencl.Device 'Intel(R) HD Graphics 5500 BroadWell
U-Processor GT2' on 'Intel Gen OCL Driver' at 0x7f4163b190e0>:
stringInput.cl:14:3: error: use of type 'double' requires
cl_khr_fp64 extension to be enabled
(options: -I /usr/lib/python2.7/dist-packages/pyopencl/cl)
(source saved as /tmp/tmpNQsE23.cl)
Given the above, I'm wondering what the best resolution would be? Would
a request to add cl_khr_fp64 to Beignet's core be possible? Or would a
better solution exist where pyopencl does not assume cl_khr_fp64 core
functionality in OpenCL 1.2+?
#!/usr/bin/env python
import numpy
import pyopencl
import pyopencl.tools
CL_CTX = pyopencl.create_some_context()
CL_DEVS = CL_CTX.devices
CL_CMDQS = [ pyopencl.CommandQueue(CL_CTX, device=d) for d in CL_DEVS ]
double_dtype, double_dtype_struct_cdecl = pyopencl.tools.match_dtype_to_c_struct(
CL_DEVS[0],
'dbl',
numpy.dtype([('double_0', numpy.float64)]))
#define pycl_offsetof(st, m) \
((size_t) ((__local char *) &(dummy.m) \
- (__local char *)&dummy ))
#if __OPENCL_C_VERSION__ < 120
#pragma OPENCL EXTENSION cl_khr_fp64: enable
#endif
#define PYOPENCL_DEFINE_CDOUBLE
typedef struct {
double double_0;
} dbl;
__kernel void get_size_and_offsets(__global size_t *result)
{
result[0] = sizeof(dbl);
__local dbl dummy;
result[1] = pycl_offsetof(dbl, double_0);
}
_______________________________________________
Beignet mailing list
[email protected]
https://lists.freedesktop.org/mailman/listinfo/beignet
