cancel
Showing results for 
Search instead for 
Did you mean: 

Archives Discussions

Raistmer
Adept II

fp64 usage in kernel

I want to have one kernel for devices that supports double precision and another one for those who lack fp64.

Will

#ifdef cl_khr_fp64

// fp64 kernel

#else

//fp32 kernel

#endif

be enough in CL file for that?

Or I should enable corresponding extension via pragma?

If yes, how I could know for what device is legal and for what is not legal to enable fp64?


0 Likes
1 Solution

#ifdef cl_khr_fp64

#pragma OPENCL EXTENSION cl_khr_fp64 : enable

#define FLOAT double

#elif defined(cl_amd_fp64)

#pragma OPENCL EXTENSION cl_amd_fp64 : enable

#define FLOAT double

#else

#define FLOAT float

#endif

kernel void FLOAT#_test(global FLOAT* a, global float* b)

{

     b[get_global_id(0)] = convert_float(a[get_global_id(0));

}

One devices that support doubles, this kernel will convert all doubles to single precision floats, and on devices that don't support doubles, it does a copy.

Message was edited by: Micah Villmow
Fixed typo from cl_khr_fp64 to the correct cl_amd_fp64.

View solution in original post

0 Likes
2 Replies

#ifdef cl_khr_fp64

#pragma OPENCL EXTENSION cl_khr_fp64 : enable

#define FLOAT double

#elif defined(cl_amd_fp64)

#pragma OPENCL EXTENSION cl_amd_fp64 : enable

#define FLOAT double

#else

#define FLOAT float

#endif

kernel void FLOAT#_test(global FLOAT* a, global float* b)

{

     b[get_global_id(0)] = convert_float(a[get_global_id(0));

}

One devices that support doubles, this kernel will convert all doubles to single precision floats, and on devices that don't support doubles, it does a copy.

Message was edited by: Micah Villmow
Fixed typo from cl_khr_fp64 to the correct cl_amd_fp64.

0 Likes

there will be quite different kernels in my case so they will be inside ifdef too, but I got the idea - define will show if device is capable or not and if it's capable extension should be enabled via pragma.

thanks!