SPIRV validation error with global variables

[maleadt]

I ran into another issue trying to switch from optimized LLVM IR to using spirv-opt:

target triple = "spir64-unknown-unknown"
@exception = private unnamed_addr constant [1 x i32] [i32 42]

This generates the following SPIRV (as disassembled by spirv-dis):

; SPIR-V
; Version: 1.0
; Generator: Khronos LLVM/SPIR-V Translator; 14
; Bound: 10
; Schema: 0
               OpCapability Addresses
               OpCapability Kernel
               OpCapability Int64
          %1 = OpExtInstImport "OpenCL.std"
               OpMemoryModel Physical64 OpenCL
               OpSource Unknown 0
               OpName %exception "exception"
               OpDecorate %exception Constant
       %uint = OpTypeInt 32 0
      %ulong = OpTypeInt 64 0
    %uint_42 = OpConstant %uint 42
    %ulong_1 = OpConstant %ulong 1
%_arr_uint_ulong_1 = OpTypeArray %uint %ulong_1
%_ptr_Function__arr_uint_ulong_1 = OpTypePointer Function %_arr_uint_ulong_1
          %7 = OpConstantComposite %_arr_uint_ulong_1 %uint_42
  %exception = OpVariable %_ptr_Function__arr_uint_ulong_1 Function %7

... which fails to validate or optimize:

error: line 16: Variables can not have a function[7] storage class outside of a function
  %exception = OpVariable %_ptr_Function__arr_uint_ulong_1 Function %7

Before, I was optimizing in LLVM and passing the SPIRV binaries generated by llvm-spirv directly to the Intel driver, containing the same kind of global variables, without any issue. It's only now that I'm switching to spirv-opt, which performs validation, that I'm running into this issue.

[MrSidims]

Thanks for the report and sorry for a long response.
The global variable in the IR snippet has a default address space, which is 0 in LLVM. For spir target it maps on OpenCL private address space which is indeed a function storage class.
So I see two issues here:

1. The GV in the IR should probably be in addrspace(1) explicitly
2. The translator should detect such incorrect IR and either try to heuristically guess the correct address space or error out.

[karolherbst]

I've hit this bug and have some easy steps to hit this behavior:

__kernel void test(int __global* in, int __global* out)
{
	*out = in[(int []){ 0,0,0,1,1, }[*in]];
}

if compiled to spirv via clang -emit-llvm -O3 -cl-std=CL3.0 -target spirv64-unknown-unknown -o /dev/stdout -c tmp.cl | llvm-spirv it generates this spirv:

; SPIR-V
; Version: 1.4
; Generator: Khronos LLVM/SPIR-V Translator; 14
; Bound: 32
; Schema: 0
               OpCapability Addresses
               OpCapability Linkage
               OpCapability Kernel
               OpCapability Int64
          %1 = OpExtInstImport "OpenCL.std"
               OpMemoryModel Physical64 OpenCL
               OpEntryPoint Kernel %27 "test" %constinit
               OpSource OpenCL_C 300000
               OpName %constinit "constinit"
               OpName %test "test"
               OpDecorate %constinit Constant
               OpDecorate %constinit Alignment 4
               OpDecorate %test LinkageAttributes "test" Export
               OpDecorate %15 FuncParamAttr NoCapture
               OpDecorate %15 FuncParamAttr NoWrite
               OpDecorate %15 Alignment 4
               OpDecorate %16 FuncParamAttr NoCapture
               OpDecorate %16 Alignment 4
               OpDecorate %28 FuncParamAttr NoCapture
               OpDecorate %28 FuncParamAttr NoWrite
               OpDecorate %28 Alignment 4
               OpDecorate %29 FuncParamAttr NoCapture
               OpDecorate %29 Alignment 4
       %uint = OpTypeInt 32 0
      %ulong = OpTypeInt 64 0
     %uint_0 = OpConstant %uint 0
     %uint_1 = OpConstant %uint 1
    %ulong_5 = OpConstant %ulong 5
    %ulong_0 = OpConstant %ulong 0
%_arr_uint_ulong_5 = OpTypeArray %uint %ulong_5
%_ptr_Function__arr_uint_ulong_5 = OpTypePointer Function %_arr_uint_ulong_5
       %void = OpTypeVoid
%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
         %13 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint
%_ptr_Function_uint = OpTypePointer Function %uint
          %8 = OpConstantComposite %_arr_uint_ulong_5 %uint_0 %uint_0 %uint_0 %uint_1 %uint_1
  %constinit = OpVariable %_ptr_Function__arr_uint_ulong_5 Function %8
       %test = OpFunction %void None %13
         %15 = OpFunctionParameter %_ptr_CrossWorkgroup_uint
         %16 = OpFunctionParameter %_ptr_CrossWorkgroup_uint
         %17 = OpLabel
         %18 = OpLoad %uint %15 Aligned 4
         %19 = OpSConvert %ulong %18
         %22 = OpInBoundsPtrAccessChain %_ptr_Function_uint %constinit %ulong_0 %19
         %23 = OpLoad %uint %22 Aligned 4
         %24 = OpSConvert %ulong %23
         %25 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %15 %24
         %26 = OpLoad %uint %25 Aligned 4
               OpStore %16 %26 Aligned 4
               OpReturn
               OpFunctionEnd
         %27 = OpFunction %void None %13
         %28 = OpFunctionParameter %_ptr_CrossWorkgroup_uint
         %29 = OpFunctionParameter %_ptr_CrossWorkgroup_uint
         %30 = OpLabel
         %31 = OpFunctionCall %void %test %28 %29
               OpReturn
               OpFunctionEnd

the intermediate llvm contains this:

@constinit = private unnamed_addr constant [5 x i32] [i32 0, i32 0, i32 0, i32 1, i32 1], align 4

[karolherbst]

given that this is constant data, it could be emitted as UniformConstant instead and everything should be fine, no? Especially as no global variable is actually used in my case.

Maybe we should change the title as this issue isn't restricted to global variables.

[karolherbst]

patch against llvm-16, not sure I like it, but the general idea is to not emit global private variables, but instead handle them once they are accessed. Not quite sure if it causes any problems though.

diff --git a/lib/SPIRV/SPIRVWriter.cpp b/lib/SPIRV/SPIRVWriter.cpp
index e8b67d5e..dbc20d02 100644
--- a/lib/SPIRV/SPIRVWriter.cpp
+++ b/lib/SPIRV/SPIRVWriter.cpp
@@ -1792,6 +1792,12 @@ LLVMToSPIRVBase::transValueWithoutDecoration(Value *V, SPIRVBasicBlock *BB,
   }
 
   if (auto GV = dyn_cast<GlobalVariable>(V)) {
+    auto AddressSpace = static_cast<SPIRAddressSpace>(GV->getAddressSpace());
+    // We can't emit private variables globally, we need to create copies of each value inside each
+    // function
+    if (AddressSpace == SPIRAS_Private && !BB)
+      return nullptr;
+
     llvm::Type *Ty = GV->getValueType();
     // Though variables with common linkage type are initialized by 0,
     // they can be represented in SPIR-V as uninitialized variables with
@@ -1852,7 +1858,6 @@ LLVMToSPIRVBase::transValueWithoutDecoration(Value *V, SPIRVBasicBlock *BB,
     }
 
     SPIRVStorageClassKind StorageClass;
-    auto AddressSpace = static_cast<SPIRAddressSpace>(GV->getAddressSpace());
     bool IsVectorCompute =
         BM->isAllowedToUseExtension(ExtensionID::SPV_INTEL_vector_compute) &&
         GV->hasAttribute(kVCMetadata::VCGlobalVariable);
@@ -1872,10 +1877,14 @@ LLVMToSPIRVBase::transValueWithoutDecoration(Value *V, SPIRVBasicBlock *BB,
       StorageClass = SPIRSPIRVAddrSpaceMap::map(AddressSpace);
     }
 
-    SPIRVType *TranslatedTy = transPointerType(Ty, GV->getAddressSpace());
+    SPIRVType *TranslatedTy = transPointerType(Ty, static_cast<unsigned int>(AddressSpace));
+
+    SPIRVBasicBlock *VarBB = nullptr;
+    if (StorageClass == StorageClassFunction)
+      VarBB = BB;
     auto BVar = static_cast<SPIRVVariable *>(
         BM->addVariable(TranslatedTy, GV->isConstant(), transLinkageType(GV),
-                        BVarInit, GV->getName().str(), StorageClass, nullptr));
+                        BVarInit, GV->getName().str(), StorageClass, VarBB));
 
     if (IsVectorCompute) {
       BVar->addDecorate(DecorationVectorComputeVariableINTEL);
@@ -3946,12 +3955,12 @@ SPIRVValue *LLVMToSPIRVBase::transIntrinsicInst(IntrinsicInst *II,
       std::vector<SPIRVValue *> Elts(TNumElts, transValue(Val, BB));
       Init = BM->addCompositeConstant(CompositeTy, Elts);
     }
-    SPIRVType *VarTy = transPointerType(AT, SPIRV::SPIRAS_Constant);
+    SPIRVType *VarTy = transPointerType(AT, SPIRV::SPIRAS_Private);
     SPIRVValue *Var = BM->addVariable(VarTy, /*isConstant*/ true,
                                       spv::internal::LinkageTypeInternal, Init,
-                                      "", StorageClassUniformConstant, nullptr);
+                                      "", StorageClassFunction, BB->getParent()->getBasicBlock(0));
     SPIRVType *SourceTy =
-        transPointerType(Val->getType(), SPIRV::SPIRAS_Constant);
+        transPointerType(Val->getType(), SPIRV::SPIRAS_Private);
     SPIRVValue *Source = BM->addUnaryInst(OpBitcast, SourceTy, Var, BB);
     SPIRVValue *Target = transValue(MSI->getRawDest(), BB);
     return BM->addCopyMemorySizedInst(Target, Source, CompositeTy->getLength(),
@@ -4544,8 +4553,8 @@ bool LLVMToSPIRVBase::transGlobalVariables() {
       continue;
     } else if (MDNode *IO = ((*I).getMetadata("io_pipe_id")))
       transGlobalIOPipeStorage(&(*I), IO);
-    else if (!transValue(&(*I), nullptr))
-      return false;
+    else
+      transValue(&(*I), nullptr);
   }
   return true;
 }
@@ -4587,6 +4596,8 @@ LLVMToSPIRVBase::collectEntryPointInterfaces(SPIRVFunction *SF, Function *F) {
   std::vector<SPIRVId> Interface;
   for (auto &GV : M->globals()) {
     const auto AS = GV.getAddressSpace();
+    if (AS == SPIRAS_Private)
+      continue;
     SPIRVModule *BM = SF->getModule();
     if (!BM->isAllowedToUseVersion(VersionNumber::SPIRV_1_4))
       if (AS != SPIRAS_Input && AS != SPIRAS_Output)
diff --git a/lib/SPIRV/libSPIRV/SPIRVModule.cpp b/lib/SPIRV/libSPIRV/SPIRVModule.cpp
index 91a2e0b8..6fdf973c 100644
--- a/lib/SPIRV/libSPIRV/SPIRVModule.cpp
+++ b/lib/SPIRV/libSPIRV/SPIRVModule.cpp
@@ -1683,7 +1683,7 @@ SPIRVInstruction *SPIRVModuleImpl::addVariable(
   SPIRVVariable *Variable = new SPIRVVariable(Type, getId(), Initializer, Name,
                                               StorageClass, BB, this);
   if (BB)
-    return addInstruction(Variable, BB);
+    return addInstruction(Variable, BB, BB->getNumInst() ? BB->getInst(0) : nullptr);
 
   add(Variable);
   if (LinkageTy != internal::LinkageTypeInternal)
-- 
2.41.0