clspv goes into infinite loop

[byzin]

Hi, clspv goes into infinite loop when I try to compile atomic float increment with atomic_cmpxchg. A global float pointer is cast to a global uint pointer, then it is passed to atomic_cmpxchg.

atomic_test.cl

__kernel void testAtomicFloatInc(__global float* result,
    const uint resolution)
{
  const uint index = get_global_id(0);
  if (index < resolution) {
    float old = *result;
    float cmp = old;
    do {
      cmp = old;
      const float value = cmp + 1.0f; // float increment
      {
        __global uint* ptr = (__global uint*)result;
        const uint c = *(const __private uint*)(&cmp);
        const uint v = *(const __private uint*)(&value);
        const uint o = atomic_cmpxchg(ptr, c, v);
        old = *(const __private float*)(&o);
      }
    } while (old != cmp);
  }
}

% clspv -O=3 --c++ --f16bit_storage --inline-entry-points --int8  atomic_test.cl -o atomic_test.spv
%4 = bitcast float addrspace(1)* %result to i32 addrspace(1)*  %4 = bitcast float addrspace(1)* %result to i32 addrspace(1)* ...

clspv prints %4 = bitcast float addrspace(1)* %result to i32 addrspace(1)* repeatedly. Could you take a look at this?

clspv: c464392367a8a769d3bd4bf5ab142fffda6d3aee

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

[alan-baker]

Fundamentally this won't work because Vulkan does not support atomics on floating point.

You could change the type of result to be a global uint pointer and use OpenCL's reinterpretation casts on the values to achieve the same effect.

[byzin]

Thank you. I'll try reinterpretation cast to float.

[fused]

I also ran into this issue:

__global const BasicData* const basic_data_header = (__global const BasicData*)((__global const char*)data + basic_data_offset);

where data is a buffer that contains BasicData as well as other structs.