clspv icon indicating copy to clipboard operation
clspv copied to clipboard

Spirv-val error of CLSPV output with program-scope "const" variable

Open hooneecho opened this issue 9 months ago • 1 comments

./clspv -cl-std=CL2.0 -inline-entry-points input.sl

const float M = 4.3f;
kernel void myKernel( global float * my_out)
{
    size_t g = get_global_id(0);
    *my_out = M;
}

https://godbolt.org/z/714de5ved

The tool spirv-val returned error for the above spirv output from CLSPV. (I was supposed to get a spv input for vulkan run, and it ended up with failing as below.)

I've done with set target as vulkan1.1 ~ 1.4. spirv-val.exe --target-env "vulkan1.1" a.spv spirv-val.exe --target-env "vulkan1.2" a.spv spirv-val.exe --target-env "vulkan1.3" a.spv spirv-val.exe --target-env "vulkan1.4" a.spv

error: line 45: [VUID-StandaloneSpirv-Uniform-06807] StorageBuffer OpVariable <id> '4[%4]' has illegal type.
From Vulkan spec:
Variables identified with the StorageBuffer storage class are used to access transparent buffer backed resources. Such variables must be typed as OpTypeStruct, or an array of this type
  %4 = OpVariable %_ptr_StorageBuffer_float StorageBuffer %float_4_30000019

https://registry.khronos.org/vulkan/specs/latest/man/html/StandaloneSpirv.html#VUID-StandaloneSpirv-Uniform-06807

Is it likely that the above command(OpVariable ~~~) was generated incorrectly? If so, What would be the reasonable fix for that?

hooneecho avatar Mar 10 '25 10:03 hooneecho

If M is declared as static the problem goes away. clspv does not seem to be handling the non-static case correctly. Presumably LLVM cannot remove the variable since it expects another file might attempt to link against it. Since linking is handled differently for clspv, probably there should be something that internalizes global variables somewhere in the flow.

alan-baker avatar Mar 25 '25 15:03 alan-baker