Spirv-val error of CLSPV output with program-scope "const" variable
./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?
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.