Forums - If I use 'half' data type in OpenCL kernel instead of 'float', can we achieve 2x performance?

Hello. First of all, apologies for the late response.

Your expectation is correct, but the answer to your question "If I use 'half' data type in OpenCL kernel instead of 'float', can we achieve 2x performance?" is "not necessarily".

"Half float" (fp16) should run at double speed compared to "float" (fp32). However for OpenCL kernels there are usually higher impact bottlenecks before you get to the level of the fp16 computations.

So for example, the overhead of submitting the workgroup might be much higher than the actual workload, in which case no benefit would be seen after optimizing the workload. Another example is that the workgroup size might not be optimal, resulting in lower "occupancy" of the shader cores. This results in the effects of memory access latency defeating any benefits of fp16 performance.

Another reason is that using "half float" could result in increase in register footprint which results in a smaller workgroup size and lower overall performance.

Finally it is possible that the thermal mitigation and clock control mechanisms built into the device are interfering with your performance measurements. You should enable performance mode in the GPU when optimizing.

Therefore to be able to see the benefits of "half float", you need to account for a lot of the points above first.

Hi,

I ran the following comparison test to do a simple memory transfer with OpenCL:

#pragma OPENCL EXTENSION cl_khr_fp16 : enable

__kernel __attribute__(( work_group_size_hint (32, 32, 1) ))
void FullPrecision(__global float* in, __global float* out)
{
        const int gtidx = get_global_id(0);
        const int gtidy = get_global_id(1);
        float4 v = vload4(0, in+gtidy*W+(gtidx<<2));
        vstore4(v, 0, out+gtidy*W+(gtidx<<2));
}

__kernel __attribute__(( work_group_size_hint (32, 32, 1) ))
void HalfPrecision(__global half* in, __global half* out)
{
        const int gtidx = get_global_id(0);
        const int gtidy = get_global_id(1);
        float8 v = vload_half8(0, in+gtidy*W+(gtidx<<3));
        vstore_half8(v, 0, out+gtidy*W+(gtidx<<3));
}

__kernel __attribute__(( work_group_size_hint (32, 32, 1) ))
void ShortPrecision(__global short* in, __global short* out)
{
        const int gtidx = get_global_id(0);
        const int gtidy = get_global_id(1);

        short8 v = vload8(0, in+gtidy*W+(gtidx<<3));
        vstore8(v, 0, out+gtidy*W+(gtidx<<3));
}


The test platform is an Adreno 420.

I followed the instruction in section 3.3.5.3.2 of document "Snapdragon(TM) OpenCL General Programming and Optimization", version  80-N8592-1 L, released on August 29, 2014, in order to enable the performance mode of the GPU.

To test the "FullPrecision" kernel, I launched 32 by 32 threads per workgroup, with a global grid of 12 by 96 workgroups, to transfer a buffer of 3072 by 3072 cl_floats. It takes 6.0ms to complete (measured using cl_event).

To test the "HalfPrecision" kernel, I launched 32 by 32 threads per workgroup, with a global grid of 24 by 96 workgroups,  to transfer a buffer of 3072 by 3072 cl_half. It takes 6.7ms to complete.

To test the "ShortPrecision" kernel, I launched 32 by 32 threads per workgroup, with a global grid of 12 by 96 workgroups, to transfer a buffer of 3072 by 3072 cl_short. It takes 3.3ms to complete.

These measurements make no sense to me, since:
1) The buffer size is large enough that the overhead time to start the kernel is not a significant factor
2) I was able to launch the same local workgroup size in all 3 tests, I did NOT use qcom-sched-rule=2 flag to complile the kernels
3) I would expect "ShortPrecision" and "HalfPrecision" to have the same speed, since they should both be transferring the same amount of data with identical layout in memory

Could you shed some light on this matter?

My apologies, the 4th and 5th paragraphs below the code snippet in the previous post should read:

To test the "FullPrecision" kernel, I launched 32 by 32 threads per workgroup, with a global grid of 24 by 96 workgroups, to transfer a buffer of 3072 by 3072 cl_floats. It takes 6.0ms to complete (measured using cl_event).

To test the "HalfPrecision" kernel, I launched 32 by 32 threads per workgroup, with a global grid of 12 by 96 workgroups,  to transfer a buffer of 3072 by 3072 cl_half. It takes 6.7ms to complete.

Seems like the problem is with vload_halfn and/or vstore_halfn on the Adreno 420. I tried rewriting the 'HalfPrecision' kernel as follows (with extension cl_khr_fp16 enabled):

__kernel __attribute__(( work_group_size_hint (32, 32, 1) ))
void HalfPrecision(__global half* in, __global half* out)
{
        const int gtidx = get_global_id(0);
        const int gtidy = get_global_id(1);
        __global half8* in8 = (__global half8*) in;
        __global half8* out8 = (__global half8*) out;

        half8 v = in8[gtidy*(W>>3)+gtidx];
        out8[gtidy*(W>>3)+gtidx] = v;

},

and now the speed of this kernel is on par with that of 'ShortPrecision'.