Using OpenCL half‑precision (fp16) for faster arithmetic on supported devices

Quick answer

If your OpenCL device reports the CL_HALF_FP capability, you can store and compute with 16‑bit floating‑point values using the half type and the native_half_* built‑in functions. This reduces memory bandwidth and can increase throughput on GPUs that natively handle fp16.

How to enable and use fp16

1. Check device support – call clGetDeviceInfo for CL_DEVICE_HALF_FP_CONFIG. A non‑zero bitmask means the device advertises fp16 support.
2. Compile with the proper language version – pass -cl-std=CL2.0 (or -cl-std=CL1.2 if you target that version) to clBuildProgram. The half type is only defined in those standards.
3. Declare half data in kernels – use half for arguments, local or global memory, and perform math with the native_half_* intrinsics to avoid promotion to 32‑bit.

Worked example

// kernel.cl
__kernel void half_mul(__global const half *a,
                       __global const half *b,
                       __global half *c,
                       const unsigned int count)
{
    int idx = get_global_id(0);
    if (idx < count) {
        // native_half_mul keeps the operation in 16‑bit precision
        c[idx] = native_half_mul(a[idx], b[idx]);
    }
}

Host side (C/C++):

cl_int err;
cl_device_id dev; /* obtained earlier */
cl_ulong half_fp_config;
err = clGetDeviceInfo(dev, CL_DEVICE_HALF_FP_CONFIG, sizeof(half_fp_config), &half_fp_config, NULL);
if (err != CL_SUCCESS || half_fp_config == 0) {
    /* fp16 not available – fall back to float */
}

const char *options = "-cl-std=CL2.0 -D cl_khr_fp16";
cl_program prog = clCreateProgramWithSource(ctx, 1, &src, NULL, &err);
clBuildProgram(prog, 1, &dev, options, NULL, NULL);
/* create buffers, set kernel args, enqueue … */

Verification steps

• Device query – after calling clGetDeviceInfo for CL_DEVICE_HALF_FP_CONFIG, inspect the returned bitmask. Common flags include CL_HALF_FP_DENORM, CL_HALF_FP_INF_NAN, and CL_HALF_FP_ROUND_TO_ZERO.
• Compile‑time check – if the build fails with an error about an undefined half type, the compiler did not see the fp16 extension (perhaps missing -cl-std=CL2.0).
• Runtime sanity – run the kernel with known inputs (e.g., 1.0 * 2.0 = 2.0) and read back the result. Compare against a reference float computation; the values should match within fp16 rounding error.
• Performance hint – time the fp16 kernel and a comparable float kernel using clGetEventProfilingInfo. A speedup indicates the device is truly executing in 16‑bit; similar times may mean the driver is promoting to float internally.

Limits and common mistakes

• Not all devices truly run fp16 in hardware – some older or integrated GPUs report CL_HALF_FP but internally convert to 32‑bit, erasing the bandwidth advantage.
• Implicit promotions – mixing half and float in the same expression causes the half to be promoted to float before the operation, negating the precision benefit and potentially adding conversion overhead.
• Rounding differences – fp16 has a much smaller dynamic range (~6.1×10⁻⁵ to 6.5×10⁴) and lower precision (about 3 decimal digits). Values outside this range become zero or infinity.
• Missing native intrinsics – using standard operators like * on half may still be compiled to float math on some platforms; prefer native_half_mul, native_half_add, etc., to stay in 16‑bit.

Practical way to check the result

After kernel execution, copy the output buffer back to host memory and compute a simple checksum (e.g., sum of all elements) using fp16 arithmetic on the host (you can emulate with a float accumulator and then cast to half). If the checksum matches the float‑based reference within the expected fp16 error margin (0.5 ulp), the kernel is likely operating in half‑precision.