Open maleadt opened 3 years ago
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:
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
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.
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
I ran into another issue trying to switch from optimized LLVM IR to using
spirv-opt
:This generates the following SPIRV (as disassembled by
spirv-dis
):... which fails to validate or optimize:
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 tospirv-opt
, which performs validation, that I'm running into this issue.