Half float in OpenCL C++

[byzin]

Hi @kpet

I have a question.
Using half float in OpenCL C++ is allowed with current clspv implementation?

A segfault happened when I built the attached half_test.cl file.

half_test.zip

% clspv -O=3 -c++ -inline-entry-points -f16bit_storage half_test.cl -o half_test.spv
[1]    12312 segmentation fault (core dumped)  clspv -O=3 -c++ -inline-entry-points -f16bit_storage half_test.cl -o

clspv version: 1fcff72

LLVM (http://llvm.org/):
  LLVM version 9.0.0svn
  Optimized build.
  Default target: x86_64-unknown-linux-gnu
  Host CPU: znver1

half_test.cl

__kernel void testLoadStoreHalf(const __global half* input1,
    const __global half* input2,
    const __global half* input3,
    const __global half* input4,
    __global half* output1,
    __global half* output2,
    __global half* output3,
    __global half* output4
    )
{
  const uint index = get_global_id(0);
  const float k = 2.0f;
  if (index == 0) {
    // scalar
    for (size_t offset = 0; offset < 2; ++offset) {
      {
        const float v = vload_half(offset, input1);
        vstore_half(k * v, offset, output1);
      }
      // vector2
      {
        const float2 v = vload_half2(offset, input2);
        vstore_half2(k * v, offset, output2);
      }
      // vector3
      {
        const float3 v = vload_half3(offset, input3);
        vstore_half3(k * v, offset, output3);
      }
      // vector4
      {
        const float4 v = vload_half4(offset, input4);
        vstore_half4(k * v, offset, output4);
      }
    }
  }
}

Thanks.

[kpet]

Thanks for the report. The short answer is that clspv doesn't support the vload/vstore built-in functions in C++ code unless the pointer is in the constant address space.

For the longer explanation: clspv's C++ language mode is essentially based on OpenCL C 2.0 + C++17. In OpenCL C 2.0, there exist only two overloads of the vload/vstore built-in functions, one taking a pointer in the generic address space (that is selected here), the other one in the constant address space. The pass dealing with OpenCL built-in functions doesn't know about the overloads that take a generic pointer which thus don't get replaced. Just adding support for these overloads to the built-in function replacement pass seems to expose bugs in the pointer bitcast replacement pass. I won't get a chance to dive into this before a few days.

[byzin]

OK then, I don't use half type for now.
Thank you for confirmation.

[rjodinchr]

This is now passing

[rjodinchr]

I made a mistake, this is still not passing

[rjodinchr]

But this is a duplicate of #1079. Let's track it there