[llvm] [SPIR-V] Fix OpName and LinkageAttributes decoration of global variables (PR #120492)
via llvm-commits
llvm-commits at lists.llvm.org
Wed Dec 18 15:16:32 PST 2024
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-spir-v
Author: Vyacheslav Levytskyy (VyacheslavLevytskyy)
<details>
<summary>Changes</summary>
This PR changes `getGlobalIdentifier()` into `getName()` value when creating a name of a global variable, and fixes generation of LinkageAttributes decoration of global variables by taking into account Private Linkage in addition to Internal.
---
Full diff: https://github.com/llvm/llvm-project/pull/120492.diff
4 Files Affected:
- (modified) llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp (+4-4)
- (modified) llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp (+6-8)
- (modified) llvm/lib/Target/SPIRV/SPIRVUtils.h (+5)
- (added) llvm/test/CodeGen/SPIRV/global-var-name-linkage.ll (+58)
``````````diff
diff --git a/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp b/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp
index 77b54219a9acc4..d2b14d6d058c92 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<UndefValue>(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<UndefValue>(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 b593b9bd1d7aab..5beb8d151c09f0 100644
--- a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
@@ -3450,7 +3450,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
@@ -3506,18 +3506,16 @@ bool SPIRVInstructionSelector::selectGlobalValue(
auto GlobalVar = cast<GlobalVariable>(GV);
assert(GlobalVar->getName() != "llvm.global.annotations");
- bool HasInit = GlobalVar->hasInitializer() &&
- !isa<UndefValue>(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 da2e24c0c9abe9..60649eac628151 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 <queue>
@@ -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<UndefValue>(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 00000000000000..93b6eb70e2d575
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/global-var-name-linkage.ll
@@ -0,0 +1,58 @@
+; 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-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 }
+
+ at G1 = private unnamed_addr addrspace(1) constant %"class.sycl::_V1::nd_item" undef, align 1
+ at g1 = addrspace(1) global i32 1, align 4
+ at g2 = internal addrspace(1) global i32 2, align 4
+ at g4 = common addrspace(1) global i32 0, align 4
+ at c1 = addrspace(2) constant [2 x i32] [i32 0, i32 1], align 4
+ at n_t = external addrspace(2) constant [256 x i32]
+ at w = addrspace(1) constant i32 0, align 4
+ at a.b = internal addrspace(2) constant [2 x i32] [i32 2, i32 3], align 4
+ at e = external addrspace(1) global i32
+ at y.z = internal addrspace(1) global i32 0, align 4
+ at 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
+}
``````````
</details>
https://github.com/llvm/llvm-project/pull/120492
More information about the llvm-commits
mailing list