OpenCL 16-bit floats

Question

How to pass 16-bit floats (cl halfs) to an OpenCL kernel in .cl file? I have in my .cl file:

__kernel void func(__global half* a, __global half* b, __global half* res){
    int i = get_global_id(0);
    res[i] = a[i] - b[i];
}

I want to pass a and b in C file and treat them as floating point values. I tried using cl_half, but it's an unsigned short, so something like:

cl_half aData[1];
aData[0] = -0.005;

won't really work, because aData[0] gets rounded to 0. How to do it properly?

Answer 1

On the GPU side, in OpenCL C, you treat the numbers as half, that part is correct. OpenCL C supports the half data type and arithmetic with it. It's better to use the half2 vector type though, as only then you get the 2x FP16 arithmetic throughput of GPUs.

On the CPU side, you treat the numbers as unsigned short (this is what cl_half is under the hood), as they are only 16 bits in size and there is no half data type in C/C++. You can't just write 32-bit FP32 literals like -0.005f into there, as they will get casted to 16-bit integer. You need proper float <-> half conversion. Then, to write numbers:

typedef unsigned short ushort;
ushort aData[1];
aData[0] = float_to_half(-0.005f);

And to read the data again:

float result = half_to_float(aData[0]);

Note that on the CPU, arithmetic with the ushort/cl_half is not supported, and you always have to first convert the numbers to float.