[llvm] [SPIRV] Unify unsized array handling for AMGCN flavoured SPIR-V (PR #175848)
Alex Voicu via llvm-commits
llvm-commits at lists.llvm.org
Mon Jan 19 11:13:23 PST 2026
https://github.com/AlexVlx updated https://github.com/llvm/llvm-project/pull/175848
>From dbd2f74c3164f4e82f3992a6a3179aef7b02a2b1 Mon Sep 17 00:00:00 2001
From: Alex Voicu <alexandru.voicu at amd.com>
Date: Sun, 11 Jan 2026 00:06:20 +0000
Subject: [PATCH 1/3] Handle `sext` and `zext` when deducing the result type
from operands.
---
llvm/lib/Target/SPIRV/SPIRVPostLegalizer.cpp | 2 ++
.../instructions/zext_sext_deduce_type.ll | 23 +++++++++++++++++++
2 files changed, 25 insertions(+)
create mode 100644 llvm/test/CodeGen/SPIRV/instructions/zext_sext_deduce_type.ll
diff --git a/llvm/lib/Target/SPIRV/SPIRVPostLegalizer.cpp b/llvm/lib/Target/SPIRV/SPIRVPostLegalizer.cpp
index 5b4ddc267c9b8..199fda0b1125d 100644
--- a/llvm/lib/Target/SPIRV/SPIRVPostLegalizer.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVPostLegalizer.cpp
@@ -202,6 +202,8 @@ static SPIRVType *deduceResultTypeFromOperands(MachineInstr *I,
switch (I->getOpcode()) {
case TargetOpcode::G_CONSTANT:
case TargetOpcode::G_ANYEXT:
+ case TargetOpcode::G_SEXT:
+ case TargetOpcode::G_ZEXT:
return deduceIntTypeFromResult(ResVReg, MIB, GR);
case TargetOpcode::G_BUILD_VECTOR:
return deduceTypeFromOperandRange(I, MIB, GR, 1, I->getNumOperands());
diff --git a/llvm/test/CodeGen/SPIRV/instructions/zext_sext_deduce_type.ll b/llvm/test/CodeGen/SPIRV/instructions/zext_sext_deduce_type.ll
new file mode 100644
index 0000000000000..8b8fb8c2b3287
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/instructions/zext_sext_deduce_type.ll
@@ -0,0 +1,23 @@
+; 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 %}
+
+; 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 %}
+
+; CHECK: %[[#I8:]] = OpTypeInt 8
+; CHECK: %[[#I64:]] = OpTypeInt 64
+; CHECK: %[[#UINT32_MAX:]] = OpConstant %[[#I64]] 4294967295
+; CHECK: %[[#SHIFT:]] = OpConstant %[[#I64]] 2097152
+; CHECK-DAG: %[[#X:]] = OpFunctionParameter %[[#I8]]
+; CHECK: %[[#Y:]] = OpFunctionParameter %[[#I64]]
+; CHECK-DAG: %[[#SEXT:]] = OpSConvert %[[#I64]] %[[#X]]
+; CHECK: %[[#AND:]] = OpBitwiseAnd %[[#I64]] %[[#SEXT]] %[[#UINT32_MAX]]
+; CHECK: %[[#]] = OpShiftRightArithmetic %[[#I64]] %[[#SHIFT]] %[[#AND]]
+
+define i64 @foo(i8 %x, i64 %y) {
+ %2 = sext i8 %x to i32
+ %3 = zext i32 %2 to i64
+ %4 = ashr i64 2097152, %3
+
+ ret i64 %4
+}
\ No newline at end of file
>From 1cecb4fd8a63ff555edceb76bd3a7563b9e75797 Mon Sep 17 00:00:00 2001
From: Alex Voicu <alexandru.voicu at amd.com>
Date: Tue, 13 Jan 2026 23:40:57 +0200
Subject: [PATCH 2/3] Unify unsized array handling for AMDGCN flavoured SPIR-V.
---
llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp | 3 ++
.../SPIRV/SPIRVLegalizeZeroSizeArrays.cpp | 3 ++
llvm/lib/Target/SPIRV/SPIRVPrepareGlobals.cpp | 54 -------------------
.../CodeGen/SPIRV/fembed-bitcode-marker.ll | 10 ++--
4 files changed, 11 insertions(+), 59 deletions(-)
diff --git a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp b/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
index 9a2b0771e4dc0..40f66f10ede82 100644
--- a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
@@ -876,6 +876,9 @@ SPIRVType *SPIRVGlobalRegistry::getOpTypeArray(uint32_t NumElems,
SPIRVType *ArrayType = nullptr;
const SPIRVSubtarget &ST =
cast<SPIRVSubtarget>(MIRBuilder.getMF().getSubtarget());
+ if (NumElems == 0 &&
+ ST.getTargetTriple().getVendor() == Triple::VendorType::AMD)
+ NumElems = UINT32_MAX;
if (NumElems != 0) {
Register NumElementsVReg =
buildConstantInt(NumElems, MIRBuilder, SpvTypeInt32, EmitIR);
diff --git a/llvm/lib/Target/SPIRV/SPIRVLegalizeZeroSizeArrays.cpp b/llvm/lib/Target/SPIRV/SPIRVLegalizeZeroSizeArrays.cpp
index 8c028ae875f9a..c7f993be24636 100644
--- a/llvm/lib/Target/SPIRV/SPIRVLegalizeZeroSizeArrays.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVLegalizeZeroSizeArrays.cpp
@@ -289,6 +289,9 @@ bool SPIRVLegalizeZeroSizeArraysImpl::runOnModule(Module &M) {
// Runtime arrays are allowed for shaders, so we don't need to do anything.
if (TM.getSubtargetImpl()->isShader())
return false;
+ // 0-sized arrays are handled differently for AMDGCN flavoured SPIRV.
+ if (M.getTargetTriple().getVendor() == Triple::VendorType::AMD)
+ return false;
// First pass: create new globals (legalizing the initializer as needed) and
// track mapping (don't erase old ones yet).
diff --git a/llvm/lib/Target/SPIRV/SPIRVPrepareGlobals.cpp b/llvm/lib/Target/SPIRV/SPIRVPrepareGlobals.cpp
index a568e725b5bf1..7a9d344c8b272 100644
--- a/llvm/lib/Target/SPIRV/SPIRVPrepareGlobals.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVPrepareGlobals.cpp
@@ -35,50 +35,6 @@ struct SPIRVPrepareGlobals : public ModulePass {
bool runOnModule(Module &M) override;
};
-bool tryExtendLLVMBitcodeMarker(GlobalVariable &Bitcode) {
- assert(Bitcode.getName() == "llvm.embedded.module");
-
- ArrayType *AT = cast<ArrayType>(Bitcode.getValueType());
- if (AT->getNumElements() != 0)
- return false;
-
- ArrayType *AT1 = ArrayType::get(AT->getElementType(), 1);
- Constant *OneEltInit = Constant::getNullValue(AT1);
- Bitcode.replaceInitializer(OneEltInit);
- 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);
- 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 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());
- NewGV->takeName(&GV);
- GV.replaceAllUsesWith(NewGV);
- GV.eraseFromParent();
-
- return true;
-}
// The backend does not support GlobalAlias. Replace aliases with their aliasees
// when possible and remove them from the module.
@@ -123,16 +79,6 @@ bool SPIRVPrepareGlobals::runOnModule(Module &M) {
Changed |= tryReplaceAliasWithAliasee(GA);
}
- const bool IsAMD = M.getTargetTriple().getVendor() == Triple::AMD;
- if (!IsAMD)
- return Changed;
-
- 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/fembed-bitcode-marker.ll b/llvm/test/CodeGen/SPIRV/fembed-bitcode-marker.ll
index 4ffdb9b7f3c7a..c39c72c73212c 100644
--- a/llvm/test/CodeGen/SPIRV/fembed-bitcode-marker.ll
+++ b/llvm/test/CodeGen/SPIRV/fembed-bitcode-marker.ll
@@ -12,11 +12,11 @@
; CHECK: OpDecorate %[[#LLVM_EMBEDDED_MODULE]] Constant
; CHECK: %[[#UCHAR:]] = OpTypeInt 8 0
; CHECK: %[[#UINT:]] = OpTypeInt 32 0
-; CHECK: %[[#ONE:]] = OpConstant %[[#UINT]] 1
-; CHECK: %[[#UCHAR_ARR_1:]] = OpTypeArray %[[#UCHAR]] %[[#ONE]]
-; CHECK: %[[#UCHAR_ARR_1_PTR:]] = OpTypePointer CrossWorkgroup %[[#UCHAR_ARR_1]]
-; CHECK: %[[#CONST_UCHAR_ARR_1:]] = OpConstantNull %[[#UCHAR_ARR_1]]
-; CHECK: %[[#LLVM_EMBEDDED_MODULE]] = OpVariable %[[#UCHAR_ARR_1_PTR]] CrossWorkgroup %[[#CONST_UCHAR_ARR_1]]
+; CHECK: %[[#UINT_MAX:]] = OpConstant %[[#UINT]] 4294967295
+; CHECK: %[[#UCHAR_ARR_UINT_MAX:]] = OpTypeArray %[[#UCHAR]] %[[#UINT_MAX]]
+; CHECK: %[[#UCHAR_ARR_UINT_MAX_PTR:]] = OpTypePointer CrossWorkgroup %[[#UCHAR_ARR_UINT_MAX]]
+; CHECK: %[[#CONST_UCHAR_ARR_UINT_MAX:]] = OpConstantNull %[[#UCHAR_ARR_UINT_MAX]]
+; CHECK: %[[#LLVM_EMBEDDED_MODULE]] = OpVariable %[[#UCHAR_ARR_UINT_MAX_PTR]] CrossWorkgroup %[[#CONST_UCHAR_ARR_UINT_MAX]]
define spir_kernel void @foo() {
entry:
>From 4a2b13f2f788243386e9e83711ca395fdb8f8722 Mon Sep 17 00:00:00 2001
From: Alex Voicu <alexandru.voicu at amd.com>
Date: Tue, 13 Jan 2026 23:54:39 +0200
Subject: [PATCH 3/3] Fix formatting.
---
llvm/lib/Target/SPIRV/SPIRVPrepareGlobals.cpp | 1 -
1 file changed, 1 deletion(-)
diff --git a/llvm/lib/Target/SPIRV/SPIRVPrepareGlobals.cpp b/llvm/lib/Target/SPIRV/SPIRVPrepareGlobals.cpp
index 7a9d344c8b272..c3cb7c5f9126d 100644
--- a/llvm/lib/Target/SPIRV/SPIRVPrepareGlobals.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVPrepareGlobals.cpp
@@ -35,7 +35,6 @@ struct SPIRVPrepareGlobals : public ModulePass {
bool runOnModule(Module &M) override;
};
-
// The backend does not support GlobalAlias. Replace aliases with their aliasees
// when possible and remove them from the module.
bool tryReplaceAliasWithAliasee(GlobalAlias &GA) {
More information about the llvm-commits
mailing list