Skip to content

[SPIR-V] Fix OpName and LinkageAttributes decoration of global variables #120492

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 3 commits into from
Jan 7, 2025

Conversation

VyacheslavLevytskyy
Copy link
Contributor

@VyacheslavLevytskyy VyacheslavLevytskyy commented Dec 18, 2024

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.

Previous implementation led to an issue with back translation of SPIR-V to LLVM IR, e.g.:

@__const.G1 = private unnamed_addr addrspace(1) constant %my_type undef
...
Fails to verify module: 'common' global may not be marked constant!
ptr addrspace(1) @"llvm-link;__const.G1"

A reproducer is included as a new test case.

@llvmbot
Copy link
Member

llvmbot commented Dec 18, 2024

@llvm/pr-subscribers-backend-spir-v

Author: Vyacheslav Levytskyy (VyacheslavLevytskyy)

Changes

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 --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 }
+
+@G1 = private unnamed_addr addrspace(1) constant %"class.sycl::_V1::nd_item" undef, 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
+}

Copy link

github-actions bot commented Dec 18, 2024

✅ With the latest revision this PR passed the undef deprecator.

@VyacheslavLevytskyy VyacheslavLevytskyy merged commit a774e7f into llvm:main Jan 7, 2025
5 of 7 checks passed
@llvm-ci
Copy link
Collaborator

llvm-ci commented Jan 7, 2025

LLVM Buildbot has detected a new failure on builder llvm-x86_64-debian-dylib running on gribozavr4 while building llvm at step 7 "test-build-unified-tree-check-llvm".

Full details are available at: https://lab.llvm.org/buildbot/#/builders/60/builds/16369

Here is the relevant piece of the build log for the reference
Step 7 (test-build-unified-tree-check-llvm) failure: test (failure)
******************** TEST 'LLVM :: ExecutionEngine/JITLink/AArch32/ELF_relocations_data.s' FAILED ********************
Exit Code: 134

Command Output (stderr):
--
RUN: at line 1: rm -rf /b/1/llvm-x86_64-debian-dylib/build/test/ExecutionEngine/JITLink/AArch32/Output/ELF_relocations_data.s.tmp && mkdir -p /b/1/llvm-x86_64-debian-dylib/build/test/ExecutionEngine/JITLink/AArch32/Output/ELF_relocations_data.s.tmp/armv6 && mkdir -p /b/1/llvm-x86_64-debian-dylib/build/test/ExecutionEngine/JITLink/AArch32/Output/ELF_relocations_data.s.tmp/armv7 && mkdir -p /b/1/llvm-x86_64-debian-dylib/build/test/ExecutionEngine/JITLink/AArch32/Output/ELF_relocations_data.s.tmp/thumbv7
+ rm -rf /b/1/llvm-x86_64-debian-dylib/build/test/ExecutionEngine/JITLink/AArch32/Output/ELF_relocations_data.s.tmp
+ mkdir -p /b/1/llvm-x86_64-debian-dylib/build/test/ExecutionEngine/JITLink/AArch32/Output/ELF_relocations_data.s.tmp/armv6
+ mkdir -p /b/1/llvm-x86_64-debian-dylib/build/test/ExecutionEngine/JITLink/AArch32/Output/ELF_relocations_data.s.tmp/armv7
+ mkdir -p /b/1/llvm-x86_64-debian-dylib/build/test/ExecutionEngine/JITLink/AArch32/Output/ELF_relocations_data.s.tmp/thumbv7
RUN: at line 2: /b/1/llvm-x86_64-debian-dylib/build/bin/llvm-mc -triple=armv6-none-linux-gnueabi -arm-add-build-attributes -filetype=obj -o /b/1/llvm-x86_64-debian-dylib/build/test/ExecutionEngine/JITLink/AArch32/Output/ELF_relocations_data.s.tmp/armv6/out.o /b/1/llvm-x86_64-debian-dylib/llvm-project/llvm/test/ExecutionEngine/JITLink/AArch32/ELF_relocations_data.s
+ /b/1/llvm-x86_64-debian-dylib/build/bin/llvm-mc -triple=armv6-none-linux-gnueabi -arm-add-build-attributes -filetype=obj -o /b/1/llvm-x86_64-debian-dylib/build/test/ExecutionEngine/JITLink/AArch32/Output/ELF_relocations_data.s.tmp/armv6/out.o /b/1/llvm-x86_64-debian-dylib/llvm-project/llvm/test/ExecutionEngine/JITLink/AArch32/ELF_relocations_data.s
RUN: at line 3: /b/1/llvm-x86_64-debian-dylib/build/bin/llvm-objdump -r /b/1/llvm-x86_64-debian-dylib/build/test/ExecutionEngine/JITLink/AArch32/Output/ELF_relocations_data.s.tmp/armv6/out.o | /b/1/llvm-x86_64-debian-dylib/build/bin/FileCheck --check-prefix=CHECK-TYPE /b/1/llvm-x86_64-debian-dylib/llvm-project/llvm/test/ExecutionEngine/JITLink/AArch32/ELF_relocations_data.s
+ /b/1/llvm-x86_64-debian-dylib/build/bin/llvm-objdump -r /b/1/llvm-x86_64-debian-dylib/build/test/ExecutionEngine/JITLink/AArch32/Output/ELF_relocations_data.s.tmp/armv6/out.o
+ /b/1/llvm-x86_64-debian-dylib/build/bin/FileCheck --check-prefix=CHECK-TYPE /b/1/llvm-x86_64-debian-dylib/llvm-project/llvm/test/ExecutionEngine/JITLink/AArch32/ELF_relocations_data.s
RUN: at line 4: /b/1/llvm-x86_64-debian-dylib/build/bin/llvm-jitlink -noexec -slab-address 0x76ff0000 -slab-allocate 10Kb -slab-page-size 4096               -abs target=0x76bbe88f -check /b/1/llvm-x86_64-debian-dylib/llvm-project/llvm/test/ExecutionEngine/JITLink/AArch32/ELF_relocations_data.s /b/1/llvm-x86_64-debian-dylib/build/test/ExecutionEngine/JITLink/AArch32/Output/ELF_relocations_data.s.tmp/armv6/out.o
+ /b/1/llvm-x86_64-debian-dylib/build/bin/llvm-jitlink -noexec -slab-address 0x76ff0000 -slab-allocate 10Kb -slab-page-size 4096 -abs target=0x76bbe88f -check /b/1/llvm-x86_64-debian-dylib/llvm-project/llvm/test/ExecutionEngine/JITLink/AArch32/ELF_relocations_data.s /b/1/llvm-x86_64-debian-dylib/build/test/ExecutionEngine/JITLink/AArch32/Output/ELF_relocations_data.s.tmp/armv6/out.o
llvm-jitlink: /b/1/llvm-x86_64-debian-dylib/llvm-project/llvm/include/llvm/ExecutionEngine/Orc/SymbolStringPool.h:285: llvm::orc::SymbolStringPool::~SymbolStringPool(): Assertion `Pool.empty() && "Dangling references at pool destruction time"' failed.
PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace.
Stack dump:
0.	Program arguments: /b/1/llvm-x86_64-debian-dylib/build/bin/llvm-jitlink -noexec -slab-address 0x76ff0000 -slab-allocate 10Kb -slab-page-size 4096 -abs target=0x76bbe88f -check /b/1/llvm-x86_64-debian-dylib/llvm-project/llvm/test/ExecutionEngine/JITLink/AArch32/ELF_relocations_data.s /b/1/llvm-x86_64-debian-dylib/build/test/ExecutionEngine/JITLink/AArch32/Output/ELF_relocations_data.s.tmp/armv6/out.o
 #0 0x00007ffa32770557 llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (/b/1/llvm-x86_64-debian-dylib/build/lib/libLLVM.so.20.0git+0xf12557)
 #1 0x00007ffa3276e00e llvm::sys::RunSignalHandlers() (/b/1/llvm-x86_64-debian-dylib/build/lib/libLLVM.so.20.0git+0xf1000e)
 #2 0x00007ffa32770c2f SignalHandler(int) Signals.cpp:0:0
 #3 0x00007ffa31848140 __restore_rt (/lib/x86_64-linux-gnu/libpthread.so.0+0x13140)
 #4 0x00007ffa3136ed51 raise (/lib/x86_64-linux-gnu/libc.so.6+0x38d51)
 #5 0x00007ffa31358537 abort (/lib/x86_64-linux-gnu/libc.so.6+0x22537)
 #6 0x00007ffa3135840f (/lib/x86_64-linux-gnu/libc.so.6+0x2240f)
 #7 0x00007ffa313676d2 (/lib/x86_64-linux-gnu/libc.so.6+0x316d2)
 #8 0x000000000043e3bc (/b/1/llvm-x86_64-debian-dylib/build/bin/llvm-jitlink+0x43e3bc)
 #9 0x00007ffa34f7a227 llvm::orc::ExecutorProcessControl::~ExecutorProcessControl() (/b/1/llvm-x86_64-debian-dylib/build/lib/libLLVM.so.20.0git+0x371c227)
#10 0x00007ffa34f7bf0f llvm::orc::SelfExecutorProcessControl::~SelfExecutorProcessControl() crtstuff.c:0:0
#11 0x00007ffa34e6c1d1 llvm::orc::ExecutionSession::~ExecutionSession() (/b/1/llvm-x86_64-debian-dylib/build/lib/libLLVM.so.20.0git+0x360e1d1)
#12 0x0000000000416180 llvm::Session::~Session() (/b/1/llvm-x86_64-debian-dylib/build/bin/llvm-jitlink+0x416180)
#13 0x00000000004213ac main (/b/1/llvm-x86_64-debian-dylib/build/bin/llvm-jitlink+0x4213ac)
#14 0x00007ffa31359d7a __libc_start_main (/lib/x86_64-linux-gnu/libc.so.6+0x23d7a)
#15 0x000000000041144a _start (/b/1/llvm-x86_64-debian-dylib/build/bin/llvm-jitlink+0x41144a)
/b/1/llvm-x86_64-debian-dylib/build/test/ExecutionEngine/JITLink/AArch32/Output/ELF_relocations_data.s.script: line 10: 1666886 Aborted                 /b/1/llvm-x86_64-debian-dylib/build/bin/llvm-jitlink -noexec -slab-address 0x76ff0000 -slab-allocate 10Kb -slab-page-size 4096 -abs target=0x76bbe88f -check /b/1/llvm-x86_64-debian-dylib/llvm-project/llvm/test/ExecutionEngine/JITLink/AArch32/ELF_relocations_data.s /b/1/llvm-x86_64-debian-dylib/build/test/ExecutionEngine/JITLink/AArch32/Output/ELF_relocations_data.s.tmp/armv6/out.o

--

********************


@VyacheslavLevytskyy
Copy link
Contributor Author

The reason for a failure in Github Actions is not related to the PR, see: #74092 for the discussion.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Projects
None yet
Development

Successfully merging this pull request may close these issues.

4 participants