[Beignet] cl_khr_fp64 on OpenCL 1.2+

[Brian Paterni]

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+?
-------------- next part --------------
A non-text attachment was scrubbed...
Name: test-double-record.py
Type: text/x-python
Size: 370 bytes
Desc: not available
URL: <https://lists.freedesktop.org/archives/beignet/attachments/20160301/771ad239/attachment.py>
-------------- next part --------------

        #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);
        }