Support for cl_khr_fp64 extension

[LukeFernandes]

After solving #303, I've encountered a problem with using doubles. clvk does not support cl_khr_fp64 (i.e. use of doubles in the kernel). My application passes down double constants using that extension...

I understand it's difficult to support lots of extensions, but I think this is a more common one. Thanks for considering.

[kpet]

Supporting the whole extension is quite a bit of work but it may be that your application only uses a small subset that's already supported. You can try adding the extension to the list there: https://github.com/kpet/clvk/blob/master/src/device.cpp#L372. It may just work. If it doesn't, feel free to attach/copy the failing kernels here, I can take a look.

[LukeFernandes]

Hi

Using a build with the extension marked as supported in device.cpp, the declaration of the extension in the kernel is still ignored by the compiler as below, and the compiler throws an error the first time double is mentioned (in the kernel void parameters list):

log is clvk-Ulpb7c/source.cl:1:26: warning: unsupported OpenCL extension 'cl_khr_fp64' - ignoring
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
                         ....
clvk-Ulpb7c/source.cl:34:98: error: use of type 'double' requires cl_khr_fp64 extension to be enabled
__kernel void gauss_3x3_div_image(__read_only image2d_t oldsrc,  const int W, const int H, const double b0, const double b1, const double c0, const double c1, const double c2, __read_only image2d_t  div, __write_only image2d_t olddst) {

[kpet]

Apologies, I'd forgotten about some recent changes to clvk and clspv, FP64 now needs to be enabled in clspv as well. Removing the line that disables FP64 in clspv (https://github.com/kpet/clvk/blob/master/src/program.cpp#L764) should do the trick.