https://github.com/jmmartinez updated https://github.com/llvm/llvm-project/pull/166952
From 8cbc3fd6403aef2cff7dbe585aa8d6762a011ef4 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Juan=20Manuel=20Martinez=20Caama=C3=B1o?= <[email protected]> Date: Fri, 7 Nov 2025 12:24:18 +0100 Subject: [PATCH 1/4] [SPIRV][SPIRVPrepareGlobals] Map AMD's dynamic LDS 0-element globals to arrays with UINT32_MAX elements In HIP, dynamic LDS globals are represented using 0-element global arrays in the __shared__ language addressspace. extern __shared__ LDS[]; These are not representable in SPIRV directly. To represent them, for AMD, we use an array with UINT32_MAX-elements. These are reverse translated to 0-element arrays later in AMD's SPIRV runtime pipeline. --- llvm/lib/Target/SPIRV/SPIRVPrepareGlobals.cpp | 27 +++++++++++++++++++ llvm/test/CodeGen/SPIRV/hip_dyn_lds.ll | 20 ++++++++++++++ 2 files changed, 47 insertions(+) create mode 100644 llvm/test/CodeGen/SPIRV/hip_dyn_lds.ll diff --git a/llvm/lib/Target/SPIRV/SPIRVPrepareGlobals.cpp b/llvm/lib/Target/SPIRV/SPIRVPrepareGlobals.cpp index c44c53129f1e0..42a9577bb2054 100644 --- a/llvm/lib/Target/SPIRV/SPIRVPrepareGlobals.cpp +++ b/llvm/lib/Target/SPIRV/SPIRVPrepareGlobals.cpp @@ -13,6 +13,7 @@ #include "SPIRV.h" +#include "llvm/ADT/STLExtras.h" #include "llvm/IR/Module.h" using namespace llvm; @@ -43,6 +44,29 @@ bool tryExtendLLVMBitcodeMarker(GlobalVariable &Bitcode) { return true; } +bool tryExtendDynamicLDSGlobal(GlobalVariable &GV) { + constexpr unsigned WorkgroupAS = 3; + const bool IsWorkgroupExternal = + GV.hasExternalLinkage() && GV.getAddressSpace() == WorkgroupAS; + if (!IsWorkgroupExternal) + return false; + + const ArrayType *AT = dyn_cast<ArrayType>(GV.getValueType()); + if (!AT || AT->getNumElements() != 0) + return false; + + constexpr auto Magic = std::numeric_limits<uint32_t>::max(); + ArrayType *NewAT = ArrayType::get(AT->getElementType(), Magic); + GlobalVariable *NewGV = new GlobalVariable( + *GV.getParent(), NewAT, GV.isConstant(), GV.getLinkage(), nullptr, "", + &GV, GV.getThreadLocalMode(), WorkgroupAS, GV.isExternallyInitialized()); + NewGV->takeName(&GV); + GV.replaceAllUsesWith(NewGV); + GV.eraseFromParent(); + + return true; +} + bool SPIRVPrepareGlobals::runOnModule(Module &M) { const bool IsAMD = M.getTargetTriple().getVendor() == Triple::AMD; if (!IsAMD) @@ -52,6 +76,9 @@ bool SPIRVPrepareGlobals::runOnModule(Module &M) { if (GlobalVariable *Bitcode = M.getNamedGlobal("llvm.embedded.module")) Changed |= tryExtendLLVMBitcodeMarker(*Bitcode); + for (GlobalVariable &GV : make_early_inc_range(M.globals())) + Changed |= tryExtendDynamicLDSGlobal(GV); + return Changed; } char SPIRVPrepareGlobals::ID = 0; diff --git a/llvm/test/CodeGen/SPIRV/hip_dyn_lds.ll b/llvm/test/CodeGen/SPIRV/hip_dyn_lds.ll new file mode 100644 index 0000000000000..f0acfdfdede9d --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/hip_dyn_lds.ll @@ -0,0 +1,20 @@ +; RUN: llc -verify-machineinstrs -mtriple=spirv64-amd-amdhsa %s -o - | FileCheck %s +; RUN: %if spirv-tools %{ llc -mtriple=spirv64-amd-amdhsa %s -o - -filetype=obj | spirv-val %} + +; CHECK: OpName %[[#LDS:]] "lds" +; CHECK: OpDecorate %[[#LDS]] LinkageAttributes "lds" Import +; CHECK: %[[#UINT:]] = OpTypeInt 32 0 +; CHECK: %[[#UINT_MAX:]] = OpConstant %[[#UINT]] 4294967295 +; CHECK: %[[#LDS_ARR_TY:]] = OpTypeArray %[[#UINT]] %[[#UINT_MAX]] +; CHECK: %[[#LDS_ARR_PTR_WG:]] = OpTypePointer Workgroup %[[#LDS_ARR_TY]] +; CHECK: %[[#LDS]] = OpVariable %[[#LDS_ARR_PTR_WG]] Workgroup + +@lds = external addrspace(3) global [0 x i32] + +define spir_kernel void @foo(ptr addrspace(4) %in, ptr addrspace(4) %out) { +entry: + %val = load i32, ptr addrspace(4) %in + %add = add i32 %val, 1 + store i32 %add, ptr addrspace(4) %out + ret void +} From 0376c3e6457061bea6ec16cb9df6789b93cac69f Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Juan=20Manuel=20Martinez=20Caama=C3=B1o?= <[email protected]> Date: Mon, 10 Nov 2025 11:08:22 +0100 Subject: [PATCH 2/4] [Review] Rename Magic->UInt32Max --- llvm/lib/Target/SPIRV/SPIRVPrepareGlobals.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/llvm/lib/Target/SPIRV/SPIRVPrepareGlobals.cpp b/llvm/lib/Target/SPIRV/SPIRVPrepareGlobals.cpp index 42a9577bb2054..2b4349e5d9e39 100644 --- a/llvm/lib/Target/SPIRV/SPIRVPrepareGlobals.cpp +++ b/llvm/lib/Target/SPIRV/SPIRVPrepareGlobals.cpp @@ -55,8 +55,8 @@ bool tryExtendDynamicLDSGlobal(GlobalVariable &GV) { if (!AT || AT->getNumElements() != 0) return false; - constexpr auto Magic = std::numeric_limits<uint32_t>::max(); - ArrayType *NewAT = ArrayType::get(AT->getElementType(), Magic); + constexpr auto UInt32Max = std::numeric_limits<uint32_t>::max(); + ArrayType *NewAT = ArrayType::get(AT->getElementType(), UInt32Max); GlobalVariable *NewGV = new GlobalVariable( *GV.getParent(), NewAT, GV.isConstant(), GV.getLinkage(), nullptr, "", &GV, GV.getThreadLocalMode(), WorkgroupAS, GV.isExternallyInitialized()); From 92a436db04c29b60e475b2fd46ec073da3803ebd Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Juan=20Manuel=20Martinez=20Caama=C3=B1o?= <[email protected]> Date: Mon, 10 Nov 2025 16:42:01 +0100 Subject: [PATCH 3/4] [Review] use storageClassToAddressSpace(SPIRV::StorageClass::Workgroup); --- llvm/lib/Target/SPIRV/SPIRVPrepareGlobals.cpp | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/llvm/lib/Target/SPIRV/SPIRVPrepareGlobals.cpp b/llvm/lib/Target/SPIRV/SPIRVPrepareGlobals.cpp index 2b4349e5d9e39..0948c75b29f5b 100644 --- a/llvm/lib/Target/SPIRV/SPIRVPrepareGlobals.cpp +++ b/llvm/lib/Target/SPIRV/SPIRVPrepareGlobals.cpp @@ -12,6 +12,7 @@ //===----------------------------------------------------------------------===// #include "SPIRV.h" +#include "SPIRVUtils.h" #include "llvm/ADT/STLExtras.h" #include "llvm/IR/Module.h" @@ -45,7 +46,8 @@ bool tryExtendLLVMBitcodeMarker(GlobalVariable &Bitcode) { } bool tryExtendDynamicLDSGlobal(GlobalVariable &GV) { - constexpr unsigned WorkgroupAS = 3; + constexpr unsigned WorkgroupAS = + storageClassToAddressSpace(SPIRV::StorageClass::Workgroup); const bool IsWorkgroupExternal = GV.hasExternalLinkage() && GV.getAddressSpace() == WorkgroupAS; if (!IsWorkgroupExternal) From 136e4ce3c1ad181cce73574e7e64e6b19951abb6 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Juan=20Manuel=20Martinez=20Caama=C3=B1o?= <[email protected]> Date: Mon, 10 Nov 2025 16:44:06 +0100 Subject: [PATCH 4/4] [Review] Add comment --- llvm/lib/Target/SPIRV/SPIRVPrepareGlobals.cpp | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/llvm/lib/Target/SPIRV/SPIRVPrepareGlobals.cpp b/llvm/lib/Target/SPIRV/SPIRVPrepareGlobals.cpp index 0948c75b29f5b..14b75d7d16a4d 100644 --- a/llvm/lib/Target/SPIRV/SPIRVPrepareGlobals.cpp +++ b/llvm/lib/Target/SPIRV/SPIRVPrepareGlobals.cpp @@ -45,6 +45,14 @@ bool tryExtendLLVMBitcodeMarker(GlobalVariable &Bitcode) { return true; } +// In HIP, dynamic LDS variables are represented using 0-element global arrays +// in the __shared__ language address-space. +// +// extern __shared__ int LDS[]; +// +// These are not representable in SPIRV directly. +// To represent them, for AMD, we use an array with UINT32_MAX-elements. +// These are reverse translated to 0-element arrays. bool tryExtendDynamicLDSGlobal(GlobalVariable &GV) { constexpr unsigned WorkgroupAS = storageClassToAddressSpace(SPIRV::StorageClass::Workgroup); _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
