OpPtrAccessChain into Private storage class sometimes generated when indexing into OpenCL module-scope constant

[dneto0]

That's only allowed when indexing into StorageBuffer or Workgroup storage class.

See test/ProgramScopeConstants/constant_array_with_function_call.cl:

constant uint b[4] = { 42, 13, 0, 15 };

uint bar(constant uint* a) { return a[get_local_id(0)]; }

void kernel __attribute__((reqd_work_group_size(4, 1, 1))) foo(global uint* a) {
  *a = bar(b);
}

Will include the following in the body of bar():

%29 = OpPtrAccessChain %_ptr_Private_uint %25 %28

---

Neil Henning says: So the issue is that:

• program-scope constant variable and constant function argument variables are the same address space in OpenCL
• we also conflate global and constant address spaces when we transform OpenCL C - > Vulkan SPIR-V (because we can't use, for instance, the UniformConstant storage class for constant variables)
• we always generate OpPtrAccessChain when a non-kernel function is indexing into a pointer as if it is an array

I think the only safe thing to do here is if you see a program-scope constant variable being passed into a sub function, then inline that sub function.

[dneto0]

cc @sheredom

[dneto0]

Another example, might have a fix in the same area:

__constant float kFirst[3] = {1.0f, 2.0f, 3.0f};
__constant float kSecond[3] = {10.0f, 11.0f, 12.0f};

kernel void foo(global float*A, int c, int i) {
 *A = c==0 ? kFirst[i] : kSecond[i];
}

Produces this kind of code:

         %33 = OpVariable %_ptr_Private__arr_float_uint_3 Private %23
         %34 = OpVariable %_ptr_Private__arr_float_uint_3 Private %27
         %35 = OpVariable %_ptr_StorageBuffer__struct_4 StorageBuffer
         %36 = OpVariable %_ptr_StorageBuffer__struct_7 StorageBuffer
         %37 = OpVariable %_ptr_StorageBuffer__struct_7 StorageBuffer
         %38 = OpFunction %void None %11
         %39 = OpLabel
         %40 = OpAccessChain %_ptr_StorageBuffer_float %35 %uint_0 %uint_0
         %41 = OpAccessChain %_ptr_StorageBuffer_uint %36 %uint_0
         %42 = OpLoad %uint %41
         %43 = OpAccessChain %_ptr_StorageBuffer_uint %37 %uint_0
         %44 = OpLoad %uint %43
         %45 = OpIEqual %bool %42 %uint_0
         %46 = OpSelect %_ptr_StorageBuffer__arr_float_uint_3 %45 %33 %34
         %47 = OpAccessChain %_ptr_StorageBuffer_float %46 %44
         %48 = OpLoad %float %47
               OpStore %40 %48
               OpReturn
               OpFunctionEnd

Note the OpSelect at %46. Its operands are pointers into Private, but its result is pointer to StorageBuffer. That's invalid. We need the initializers for %33 and %34 but even with VariablePointers we can't select between two different pointer-to-Private values.

[dneto0]

We get pretty much the identical SPIR-V when we pull the __constant values inside the function as private variables:

kernel void foo(global float *A, int c, int i) {
  float kFirst[3] = {1.0f, 2.0f, 3.0f};
  float kSecond[3] = {10.0f, 11.0f, 12.0f};
  *A = c == 0 ? kFirst[i] : kSecond[i];
}

Results in:

         %33 = OpVariable %_ptr_Private__arr_float_uint_3 Private %23
         %34 = OpVariable %_ptr_Private__arr_float_uint_3 Private %27
         %35 = OpVariable %_ptr_StorageBuffer__struct_4 StorageBuffer
         %36 = OpVariable %_ptr_StorageBuffer__struct_7 StorageBuffer
         %37 = OpVariable %_ptr_StorageBuffer__struct_7 StorageBuffer
         %38 = OpFunction %void None %11
         %39 = OpLabel
         %40 = OpAccessChain %_ptr_StorageBuffer_float %35 %uint_0 %uint_0
         %41 = OpAccessChain %_ptr_StorageBuffer_uint %36 %uint_0
         %42 = OpLoad %uint %41
         %43 = OpAccessChain %_ptr_StorageBuffer_uint %37 %uint_0
         %44 = OpLoad %uint %43
         %45 = OpIEqual %bool %42 %uint_0
         %46 = OpSelect %_ptr_StorageBuffer__arr_float_uint_3 %45 %33 %34
         %47 = OpAccessChain %_ptr_StorageBuffer_float %46 %44
         %48 = OpLoad %float %47

For this code, it seems the compiler has gone too far in inst-combine. It has changed a selection of the two loaded values into a selection of the two pointers, and then a load from that pointer. Perhaps a small workaround is to reverse that decision, at least for this case.

[dneto0]

A workaround for the kFirst kSecond example is to wrap the accesses in functions, like this:

float first(int i) {
  float kFirst[3] = {1.0f, 2.0f, 3.0f};
  return kFirst[i];
}

float second(int i) {
  float kSecond[3] = {10.0f, 11.0f, 12.0f};
  return kSecond[i];
}

kernel void foo(global float *A, int c, int i) {
  *A = c == 0 ? first(i) : second(i);
}

[dneto0]

Another workaround for the kFirst kSecond example is to make the decision an index:

kernel void foo(global float *A, int c, int i) {
  float kLookup[2][3] = {{1.0f, 2.0f, 3.0f}, {10.0f, 11.0f, 12.0f}};
  *A = kLookup[c != 0][i];
}

[dneto0]

I've forked the issue with OpSelect between pointers into #71. This issue remains open for the call-to-function with pointer to private.

[PlasmaPower]

I think I ran into this while attempting to reproduce #728 . Kernel: https://gist.github.com/PlasmaPower/1c7e9efba6fa36ca6ef7156e335825bb

OpPtrAccessChain is not supported for this storage class
UNREACHABLE executed at ../../lib/SPIRVProducerPass.cpp:3907!

[PlasmaPower]

I've reduced my issue. Interestingly, it doesn't seem to be caused by indexing into an constant array. There are a lot of necessary parts to this, but it does still have a constant, albeit not a constant array.

__constant uint reduce_mask_26 = (1 << 26) - 1;

void curve25519_add_reduce(uint out[10]) {
    out[8] = reduce_mask_26;
}

int ed25519_verify(const unsigned char *x, size_t len) {
    size_t differentbits = 0;
    while (len--)
        differentbits |= *x++;
    return differentbits;
}

void ge25519_unpack_vartime() {
    unsigned char check[32];
    uint t[10];
    if (!ed25519_verify(check, 32)) {
        curve25519_add_reduce(t);
    }
}

__kernel void generate_pubkey () {
    ge25519_unpack_vartime();
}

[rjodinchr]

I confirm that this issue is still relevant

[rjodinchr]

I am not able to reproduce the issue using clspv ToT. Should we close this issue?