Double precision complex numbers

[tnorth]

Hello,

First thanks for that great piece of software.
I am trying to use double precision complex with PyOpenCL (complex128) on an ATI platform. I guess I am not doing it properly, because I get an error message (see below). Or is it an issue with that OpenCL extension cl_khr_fp64 ?
Adding a preamble like in http://documen.tician.de/pyopencl/array.html#complex-numbers leads to the same results.

On the CPU (intel i7), the following will work:

 norm2 = ElementwiseKernel(ctx,
            "cdouble_t *x, cdouble_t *z",
            "z[i] = cdouble_mul(x[i], cdouble_conj(x[i]))",
            "norm2")

Also, it seems that many test of pyopencl/test/ fail for the same issue.
I am using the last GIT of PyOpenCL on linux 3.2.7 x86_64, gcc version 4.6.2, ati catalyst 12.1, OpenCL SDK 2.6.
Any hint ?
Thanks,

norm2 = ElementwiseKernel(ctx,
    "cdouble_t *x, cdouble_t *z",
    "z[i] = cdouble_mul(x[i], cdouble_conj(x[i]))",
    "norm2")
## -- End pasted text --
---------------------------------------------------------------------------
RuntimeError                              Traceback (most recent call last)
/home/tnorth/Documents/phd/conferences/UP2012/simus/ssfs_20120209_beta3/<ipython-input-93-22eceb07b4a1> in <module>()
      2     "cdouble_t *x, cdouble_t *z",
      3     "z[i] = cdouble_mul(x[i], cdouble_conj(x[i]))",
----> 4     "norm2")

/usr/lib64/python2.7/site-packages/pyopencl-2011.2-py2.7-linux-x86_64.egg/pyopencl/elementwise.pyc in __init__(self, context, arguments, operation, name, options, **kwargs)                                                              
    141             context, arguments, operation,
    142             name=name, options=options,
--> 143             **kwargs)
    144                                                                                                              
    145         if not [i for i, arg in enumerate(self.arguments)

/usr/lib64/python2.7/site-packages/pyopencl-2011.2-py2.7-linux-x86_64.egg/pyopencl/elementwise.pyc in get_elwise_kernel_and_types(context, arguments, operation, name, options, preamble, **kwargs)                                       
    102     prg = get_elwise_program(
    103         context, parsed_args, operation,
--> 104         name=name, options=options, preamble=preamble, **kwargs)
    105                                                                                                              
    106     scalar_arg_dtypes = []

/usr/lib64/python2.7/site-packages/pyopencl-2011.2-py2.7-linux-x86_64.egg/pyopencl/elementwise.pyc in get_elwise_program(context, arguments, operation, name, options, preamble, loop_prep, after_loop)                                   
     67             })
     68 
---> 69     return Program(context, source).build(options)
     70 
     71 

/usr/lib64/python2.7/site-packages/pyopencl-2011.2-py2.7-linux-x86_64.egg/pyopencl/__init__.pyc in build(self, options, devices, cache_dir)                                                                                               
    122             self._prg = create_built_program_from_source_cached(
    123                     self._context, self._source, options, devices,
--> 124                     cache_dir=cache_dir)
    125                                                                                                              
    126         return self

/usr/lib64/python2.7/site-packages/pyopencl-2011.2-py2.7-linux-x86_64.egg/pyopencl/cache.pyc in create_built_program_from_source_cached(ctx, src, options, devices, cache_dir)                                                            
    457         if cache_dir != False:
    458             prg, already_built = _create_built_program_from_source_cached(
--> 459                     ctx, src, options, devices, cache_dir)
    460         else:                                                                                                
    461             prg = _cl._Program(ctx, src)

/usr/lib64/python2.7/site-packages/pyopencl-2011.2-py2.7-linux-x86_64.egg/pyopencl/cache.pyc in _create_built_program_from_source_cached(ctx, src, options, devices, cache_dir)                                                           
    381 
    382         prg = _cl._Program(ctx, src)
--> 383         prg.build(options, [devices[i] for i in to_be_built_indices])
    384 
    385         prg_devs = prg.get_info(_cl.program_info.DEVICES)

/usr/lib64/python2.7/site-packages/pyopencl-2011.2-py2.7-linux-x86_64.egg/pyopencl/__init__.pyc in program_build(self, options, devices)                                                                                                  
    375             # Python 3.2 outputs the whole list of currently active exceptions

    376             # This serves to remove one (redundant) level from that nesting.

--> 377             raise err
    378 
    379         message = (75*"="+"\n").join(

RuntimeError: clBuildProgram failed: build program failure - 

Build on <pyopencl.Device 'Juniper' on 'AMD Accelerated Parallel Processing' at 0x1647f20>:

/tmp/OCLLZXKXc.cl(2): error: can't enable all OpenCL extensions or
          unrecognized OpenCL extension
          #pragma OPENCL EXTENSION cl_khr_fp64: enable
                                                ^

/usr/lib64/python2.7/site-packages/pyopencl-2011.2-py2.7-linux-x86_64.egg/pyopencl/../include/pyopencl/pyopencl-complex.h(223): error: 
          identifier "double2" is undefined
  PYOPENCL_DECLARE_COMPLEX_TYPE(double, DBL);
  ^

[tnorth]

Hmm it seems that this card does not support the cl_amd_pf64 extension according to ATI's guide....
Is there a workaround though ?

[inducer]

Nope--"no double precision" is not something PyOpenCL is designed to work around.

HTH,
Andreas