google
Half float in OpenCL C++
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.
% 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.
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.
