diff --git a/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp b/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp index 77b54219a9acc..d2b14d6d058c9 100644 --- a/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp +++ b/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp @@ -1841,20 +1841,20 @@ void SPIRVEmitIntrinsics::processGlobalValue(GlobalVariable &GV, // Skip special artifical variable llvm.global.annotations. if (GV.getName() == "llvm.global.annotations") return; - if (GV.hasInitializer() && !isa(GV.getInitializer())) { + Constant *Init = nullptr; + if (hasInitializer(&GV)) { // Deduce element type and store results in Global Registry. // Result is ignored, because TypedPointerType is not supported // by llvm IR general logic. deduceElementTypeHelper(&GV, false); - Constant *Init = GV.getInitializer(); + Init = GV.getInitializer(); Type *Ty = isAggrConstForceInt32(Init) ? B.getInt32Ty() : Init->getType(); Constant *Const = isAggrConstForceInt32(Init) ? B.getInt32(1) : Init; auto *InitInst = B.CreateIntrinsic(Intrinsic::spv_init_global, {GV.getType(), Ty}, {&GV, Const}); InitInst->setArgOperand(1, Init); } - if ((!GV.hasInitializer() || isa(GV.getInitializer())) && - GV.getNumUses() == 0) + if (!Init && GV.getNumUses() == 0) B.CreateIntrinsic(Intrinsic::spv_unref_global, GV.getType(), &GV); } diff --git a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp index db970275ebd5b..237f71a1b70e5 100644 --- a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp +++ b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp @@ -3478,7 +3478,7 @@ bool SPIRVInstructionSelector::selectGlobalValue( ID = UnnamedGlobalIDs.size(); GlobalIdent = "__unnamed_" + Twine(ID).str(); } else { - GlobalIdent = GV->getGlobalIdentifier(); + GlobalIdent = GV->getName(); } // Behaviour of functions as operands depends on availability of the @@ -3541,18 +3541,16 @@ bool SPIRVInstructionSelector::selectGlobalValue( auto GlobalVar = cast(GV); assert(GlobalVar->getName() != "llvm.global.annotations"); - bool HasInit = GlobalVar->hasInitializer() && - !isa(GlobalVar->getInitializer()); - // Skip empty declaration for GVs with initilaizers till we get the decl with + // Skip empty declaration for GVs with initializers till we get the decl with // passed initializer. - if (HasInit && !Init) + if (hasInitializer(GlobalVar) && !Init) return true; - bool HasLnkTy = GV->getLinkage() != GlobalValue::InternalLinkage; + bool HasLnkTy = !GV->hasInternalLinkage() && !GV->hasPrivateLinkage(); SPIRV::LinkageType::LinkageType LnkType = - (GV->isDeclaration() || GV->hasAvailableExternallyLinkage()) + GV->isDeclarationForLinker() ? SPIRV::LinkageType::Import - : (GV->getLinkage() == GlobalValue::LinkOnceODRLinkage && + : (GV->hasLinkOnceODRLinkage() && STI.canUseExtension(SPIRV::Extension::SPV_KHR_linkonce_odr) ? SPIRV::LinkageType::LinkOnceODR : SPIRV::LinkageType::Export); diff --git a/llvm/lib/Target/SPIRV/SPIRVUtils.h b/llvm/lib/Target/SPIRV/SPIRVUtils.h index da2e24c0c9abe..60649eac62815 100644 --- a/llvm/lib/Target/SPIRV/SPIRVUtils.h +++ b/llvm/lib/Target/SPIRV/SPIRVUtils.h @@ -17,6 +17,7 @@ #include "llvm/Analysis/LoopInfo.h" #include "llvm/CodeGen/MachineBasicBlock.h" #include "llvm/IR/Dominators.h" +#include "llvm/IR/GlobalVariable.h" #include "llvm/IR/IRBuilder.h" #include "llvm/IR/TypedPointerType.h" #include @@ -236,6 +237,10 @@ Type *parseBasicTypeName(StringRef &TypeName, LLVMContext &Ctx); // Returns true if the function was changed. bool sortBlocks(Function &F); +inline bool hasInitializer(const GlobalVariable *GV) { + return GV->hasInitializer() && !isa(GV->getInitializer()); +} + // True if this is an instance of TypedPointerType. inline bool isTypedPointerTy(const Type *T) { return T && T->getTypeID() == Type::TypedPointerTyID; diff --git a/llvm/test/CodeGen/SPIRV/global-var-name-linkage.ll b/llvm/test/CodeGen/SPIRV/global-var-name-linkage.ll new file mode 100644 index 0000000000000..4501819ce4940 --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/global-var-name-linkage.ll @@ -0,0 +1,59 @@ +; Check names and decoration of global variables. + +; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s +; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} + +; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s +; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %} + +; CHECK-DAG: OpName %[[#id18:]] "G1" +; CHECK-DAG: OpName %[[#id22:]] "g1" +; CHECK-DAG: OpName %[[#id23:]] "g2" +; CHECK-DAG: OpName %[[#id27:]] "g4" +; CHECK-DAG: OpName %[[#id30:]] "c1" +; CHECK-DAG: OpName %[[#id31:]] "n_t" +; CHECK-DAG: OpName %[[#id32:]] "w" +; CHECK-DAG: OpName %[[#id34:]] "a.b" +; CHECK-DAG: OpName %[[#id35:]] "e" +; CHECK-DAG: OpName %[[#id36:]] "y.z" +; CHECK-DAG: OpName %[[#id38:]] "x" + +; CHECK-NOT: OpDecorate %[[#id18]] LinkageAttributes +; CHECK-DAG: OpDecorate %[[#id18]] Constant +; CHECK-DAG: OpDecorate %[[#id22]] Alignment 4 +; CHECK-DAG: OpDecorate %[[#id22]] LinkageAttributes "g1" Export +; CHECK-DAG: OpDecorate %[[#id23]] Alignment 4 +; CHECK-DAG: OpDecorate %[[#id27]] Alignment 4 +; CHECK-DAG: OpDecorate %[[#id27]] LinkageAttributes "g4" Export +; CHECK-DAG: OpDecorate %[[#id30]] Constant +; CHECK-DAG: OpDecorate %[[#id30]] Alignment 4 +; CHECK-DAG: OpDecorate %[[#id30]] LinkageAttributes "c1" Export +; CHECK-DAG: OpDecorate %[[#id31]] Constant +; CHECK-DAG: OpDecorate %[[#id31]] LinkageAttributes "n_t" Import +; CHECK-DAG: OpDecorate %[[#id32]] Constant +; CHECK-DAG: OpDecorate %[[#id32]] Alignment 4 +; CHECK-DAG: OpDecorate %[[#id32]] LinkageAttributes "w" Export +; CHECK-DAG: OpDecorate %[[#id34]] Constant +; CHECK-DAG: OpDecorate %[[#id34]] Alignment 4 +; CHECK-DAG: OpDecorate %[[#id35]] LinkageAttributes "e" Import +; CHECK-DAG: OpDecorate %[[#id36]] Alignment 4 +; CHECK-DAG: OpDecorate %[[#id38]] Constant +; CHECK-DAG: OpDecorate %[[#id38]] Alignment 4 + +%"class.sycl::_V1::nd_item" = type { i8 } + +@G1 = private unnamed_addr addrspace(1) constant %"class.sycl::_V1::nd_item" poison, align 1 +@g1 = addrspace(1) global i32 1, align 4 +@g2 = internal addrspace(1) global i32 2, align 4 +@g4 = common addrspace(1) global i32 0, align 4 +@c1 = addrspace(2) constant [2 x i32] [i32 0, i32 1], align 4 +@n_t = external addrspace(2) constant [256 x i32] +@w = addrspace(1) constant i32 0, align 4 +@a.b = internal addrspace(2) constant [2 x i32] [i32 2, i32 3], align 4 +@e = external addrspace(1) global i32 +@y.z = internal addrspace(1) global i32 0, align 4 +@x = internal addrspace(2) constant float 1.000000e+00, align 4 + +define internal spir_func void @foo(ptr addrspace(4) align 1 %arg) { + ret void +}