Built-in variables used through a pointer (bitcast/access chain) are never declared

[maleadt]

Built-in variables are only declared in the generated OpenCL C when they are loaded directly. If a built-in's pointer is used any other way, e.g. bitcast and indexed to read one component, the C code refers to the SPIR-V variable name, which is never declared, and the program fails to build.

The LLVM SPIR-V back end emits this for LLVM IR that loads a single component of a 3-D built-in through a GEP:

%ptr = getelementptr inbounds <3 x i64>, ptr addrspace(7) @__spirv_BuiltInLocalInvocationId, i64 0, i64 %idx
%val = load i64, ptr addrspace(7) %ptr

SPIR-V (passes spirv-val, runs correctly on PoCL):

               OpDecorate %__spirv_BuiltInLocalInvocationId BuiltIn LocalInvocationId
%__spirv_BuiltInLocalInvocationId = OpVariable %_ptr_Input_v3ulong Input
         %31 = OpBitcast %_ptr_Input_uchar %__spirv_BuiltInLocalInvocationId
         %32 = OpInBoundsPtrAccessChain %_ptr_Input_uchar %31 %ulong_8
         %33 = OpBitcast %_ptr_Input_uchar %__spirv_BuiltInLocalInvocationId
         %34 = OpInBoundsPtrAccessChain %_ptr_Input_uchar %33 %ulong_16
         %39 = OpBitcast %_ptr_Input_ulong %__spirv_BuiltInLocalInvocationId
         %40 = OpLoad %ulong %39 Aligned 4
         %41 = OpUConvert %uint %40
         %42 = OpBitcast %_ptr_Input_ulong %32
         %43 = OpLoad %ulong %42 Aligned 4

Full module

               OpCapability Kernel
               OpCapability Addresses
               OpCapability Int64
               OpCapability Int8
               OpCapability SubgroupDispatch
          %1 = OpExtInstImport "OpenCL.std"
               OpMemoryModel Physical64 OpenCL
               OpEntryPoint Kernel %_Z10gep_kernel13CLDeviceArrayI5Int32Li1ELi1EE "_Z10gep_kernel13CLDeviceArrayI5Int32Li1ELi1EE" %__spirv_BuiltInLocalInvocationId
               OpExecutionMode %_Z10gep_kernel13CLDeviceArrayI5Int32Li1ELi1EE SubgroupSize 32
               OpExecutionMode %_Z10gep_kernel13CLDeviceArrayI5Int32Li1ELi1EE ContractionOff
               OpSource OpenCL_C 200000
               OpName %state_ptr "state_ptr"
               OpName %a__CLDeviceArray "a::CLDeviceArray"
               OpName %_Z10gep_kernel13CLDeviceArrayI5Int32Li1ELi1EE "_Z10gep_kernel13CLDeviceArrayI5Int32Li1ELi1EE"
               OpName %__spirv_BuiltInLocalInvocationId "__spirv_BuiltInLocalInvocationId"
               OpName %conversion "conversion"
               OpDecorate %state_ptr FuncParamAttr ByVal
               OpDecorate %a__CLDeviceArray FuncParamAttr ByVal
               OpDecorate %__spirv_BuiltInLocalInvocationId BuiltIn LocalInvocationId
      %ulong = OpTypeInt 64 0
       %uint = OpTypeInt 32 0
  %_struct_5 = OpTypeStruct %uint %ulong %ulong
  %_struct_6 = OpTypeStruct %_struct_5
%_ptr_Function__struct_6 = OpTypePointer Function %_struct_6
      %uchar = OpTypeInt 8 0
%_ptr_CrossWorkgroup_uchar = OpTypePointer CrossWorkgroup %uchar
     %uint_1 = OpConstant %uint 1
%_arr_ulong_uint_1 = OpTypeArray %ulong %uint_1
 %_struct_12 = OpTypeStruct %_ptr_CrossWorkgroup_uchar %ulong %_arr_ulong_uint_1 %ulong
 %_struct_13 = OpTypeStruct %_struct_12
%_ptr_Function__struct_13 = OpTypePointer Function %_struct_13
       %void = OpTypeVoid
         %16 = OpTypeFunction %void %_ptr_Function__struct_6 %_ptr_Function__struct_13
%_ptr_Input_ulong = OpTypePointer Input %ulong
%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
%_ptr_Input_uchar = OpTypePointer Input %uchar
    %v3ulong = OpTypeVector %ulong 3
%_ptr_Input_v3ulong = OpTypePointer Input %v3ulong
    %ulong_2 = OpConstant %ulong 2
   %ulong_16 = OpConstant %ulong 16
    %ulong_1 = OpConstant %ulong 1
    %ulong_8 = OpConstant %ulong 8
%_ptr_Function__ptr_CrossWorkgroup_uint = OpTypePointer Function %_ptr_CrossWorkgroup_uint
%__spirv_BuiltInLocalInvocationId = OpVariable %_ptr_Input_v3ulong Input
%_Z10gep_kernel13CLDeviceArrayI5Int32Li1ELi1EE = OpFunction %void None %16
  %state_ptr = OpFunctionParameter %_ptr_Function__struct_6
%a__CLDeviceArray = OpFunctionParameter %_ptr_Function__struct_13
 %conversion = OpLabel
         %31 = OpBitcast %_ptr_Input_uchar %__spirv_BuiltInLocalInvocationId
         %32 = OpInBoundsPtrAccessChain %_ptr_Input_uchar %31 %ulong_8
         %33 = OpBitcast %_ptr_Input_uchar %__spirv_BuiltInLocalInvocationId
         %34 = OpInBoundsPtrAccessChain %_ptr_Input_uchar %33 %ulong_16
         %35 = OpBitcast %_ptr_Function__ptr_CrossWorkgroup_uint %a__CLDeviceArray
         %36 = OpLoad %_ptr_CrossWorkgroup_uint %35 Aligned 8
         %37 = OpBitcast %_ptr_CrossWorkgroup_uchar %36
         %38 = OpBitcast %_ptr_CrossWorkgroup_uchar %36
         %39 = OpBitcast %_ptr_Input_ulong %__spirv_BuiltInLocalInvocationId
         %40 = OpLoad %ulong %39 Aligned 4
         %41 = OpUConvert %uint %40
               OpStore %36 %41 Aligned 4
         %42 = OpBitcast %_ptr_Input_ulong %32
         %43 = OpLoad %ulong %42 Aligned 4
         %44 = OpUConvert %uint %43
         %45 = OpBitcast %_ptr_CrossWorkgroup_uint %37
         %46 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %45 %ulong_1
               OpStore %46 %44 Aligned 4
         %47 = OpBitcast %_ptr_Input_ulong %34
         %48 = OpLoad %ulong %47 Aligned 4
         %49 = OpUConvert %uint %48
         %50 = OpBitcast %_ptr_CrossWorkgroup_uint %38
         %51 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %50 %ulong_2
               OpStore %51 %49 Aligned 4
               OpReturn
               OpFunctionEnd

Generated OpenCL C (spirv2clc 0.2.1, spirv2clc_jll v0.2.1+0), rejected by NVIDIA's OpenCL compiler:

  ma8 * v31 = ((ma8 *)__spirv_BuiltInLocalInvocationId);
  ma8 * v32 = &v31[(long)(((ulong)8))];
  ma8 * v33 = ((ma8 *)__spirv_BuiltInLocalInvocationId);
  ma8 * v34 = &v33[(long)(((ulong)16))];
  ...
  ma3 * v39 = ((ma3 *)__spirv_BuiltInLocalInvocationId);
  ulong v40 = *((ma3a4 *)v39);

<kernel>:37:23: error: use of undeclared identifier '__spirv_BuiltInLocalInvocationId'
<kernel>:39:23: error: use of undeclared identifier '__spirv_BuiltInLocalInvocationId'
<kernel>:45:23: error: use of undeclared identifier '__spirv_BuiltInLocalInvocationId'

Expected: the built-in materialized as a private variable (e.g. ulong3 __spirv_BuiltInLocalInvocationId = (ulong3)(get_local_id(0), get_local_id(1), get_local_id(2));, or per component), so that pointer uses of it work like those of any other variable.

Context: SPIRVIntrinsics loaded single components this way while working around an InstCombine fold that produces <24 x i8> loads for truncated built-ins; it now calls __spirv_BuiltIn* functions instead, so this isn't blocking.