[llvm] [SPIRV] Unify unsized array handling for AMGCN flavoured SPIR-V (PR #175848)
Alex Voicu via llvm-commits
llvm-commits at lists.llvm.org
Tue Jan 20 03:24:11 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/6] 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/6] 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/6] 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) {
>From ce153ae2b3372f0f59e998c0cce2ad0f5afac273 Mon Sep 17 00:00:00 2001
From: Alex Voicu <alexandru.voicu at amd.com>
Date: Tue, 20 Jan 2026 00:53:15 +0000
Subject: [PATCH 4/6] Use `UINT64_MAX`, as it's guaranteed to be illegal /
invalid/
---
llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp | 15 ++++++++++++---
llvm/test/CodeGen/SPIRV/fembed-bitcode-marker.ll | 12 ++++++------
llvm/test/CodeGen/SPIRV/hip_dyn_lds.ll | 5 +++--
3 files changed, 21 insertions(+), 11 deletions(-)
diff --git a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp b/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
index ad28bc0625fc3..5200f1ead9a13 100644
--- a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
@@ -876,9 +876,6 @@ 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);
@@ -888,6 +885,18 @@ SPIRVType *SPIRVGlobalRegistry::getOpTypeArray(uint32_t NumElems,
.addUse(getSPIRVTypeID(ElemType))
.addUse(NumElementsVReg);
});
+ } else if (ST.getTargetTriple().getVendor() == Triple::VendorType::AMD) {
+ // We set the array size to the token UINT64_MAX value, which is generally
+ // illegal (the maximum legal size is 61-bits) for the foreseeable future.
+ SPIRVType *SpvTypeInt64 = getOrCreateSPIRVIntegerType(64, MIRBuilder);
+ Register NumElementsVReg =
+ buildConstantInt(UINT64_MAX, MIRBuilder, SpvTypeInt64, EmitIR);
+ ArrayType = createOpType(MIRBuilder, [&](MachineIRBuilder &MIRBuilder) {
+ return MIRBuilder.buildInstr(SPIRV::OpTypeArray)
+ .addDef(createTypeVReg(MIRBuilder))
+ .addUse(getSPIRVTypeID(ElemType))
+ .addUse(NumElementsVReg);
+ });
} else {
if (!ST.isShader()) {
llvm::reportFatalUsageError(
diff --git a/llvm/test/CodeGen/SPIRV/fembed-bitcode-marker.ll b/llvm/test/CodeGen/SPIRV/fembed-bitcode-marker.ll
index c39c72c73212c..b17464199c720 100644
--- a/llvm/test/CodeGen/SPIRV/fembed-bitcode-marker.ll
+++ b/llvm/test/CodeGen/SPIRV/fembed-bitcode-marker.ll
@@ -11,12 +11,12 @@
; CHECK: OpName %[[#LLVM_EMBEDDED_MODULE:]] "llvm.embedded.module"
; CHECK: OpDecorate %[[#LLVM_EMBEDDED_MODULE]] Constant
; CHECK: %[[#UCHAR:]] = OpTypeInt 8 0
-; CHECK: %[[#UINT:]] = OpTypeInt 32 0
-; 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]]
+; CHECK: %[[#UINT64:]] = OpTypeInt 64 0
+; CHECK: %[[#UINT64_MAX:]] = OpConstant %[[#UINT64]] 18446744073709551615
+; CHECK: %[[#UCHAR_ARR_UINT64_MAX:]] = OpTypeArray %[[#UCHAR]] %[[#UINT64_MAX]]
+; CHECK: %[[#UCHAR_ARR_UINT64_MAX_PTR:]] = OpTypePointer CrossWorkgroup %[[#UCHAR_ARR_UINT64_MAX]]
+; CHECK: %[[#CONST_UCHAR_ARR_UINT64_MAX:]] = OpConstantNull %[[#UCHAR_ARR_UINT64_MAX]]
+; CHECK: %[[#LLVM_EMBEDDED_MODULE]] = OpVariable %[[#UCHAR_ARR_UINT64_MAX_PTR]] CrossWorkgroup %[[#CONST_UCHAR_ARR_UINT64_MAX]]
define spir_kernel void @foo() {
entry:
diff --git a/llvm/test/CodeGen/SPIRV/hip_dyn_lds.ll b/llvm/test/CodeGen/SPIRV/hip_dyn_lds.ll
index f0acfdfdede9d..531d5817ee86c 100644
--- a/llvm/test/CodeGen/SPIRV/hip_dyn_lds.ll
+++ b/llvm/test/CodeGen/SPIRV/hip_dyn_lds.ll
@@ -4,8 +4,9 @@
; 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: %[[#UINT64:]] = OpTypeInt 64 0
+; CHECK: %[[#UINT64_MAX:]] = OpConstant %[[#UINT64]] 18446744073709551615
+; CHECK: %[[#LDS_ARR_TY:]] = OpTypeArray %[[#UINT]] %[[#UINT64_MAX]]
; CHECK: %[[#LDS_ARR_PTR_WG:]] = OpTypePointer Workgroup %[[#LDS_ARR_TY]]
; CHECK: %[[#LDS]] = OpVariable %[[#LDS_ARR_PTR_WG]] Workgroup
>From e3ea50015011b684f7f4926bfd02212100120e99 Mon Sep 17 00:00:00 2001
From: Alex Voicu <alexandru.voicu at amd.com>
Date: Tue, 20 Jan 2026 10:57:53 +0000
Subject: [PATCH 5/6] Adopt review suggestion.
---
llvm/lib/Target/SPIRV/SPIRVLegalizeZeroSizeArrays.cpp | 4 ----
1 file changed, 4 deletions(-)
diff --git a/llvm/lib/Target/SPIRV/SPIRVLegalizeZeroSizeArrays.cpp b/llvm/lib/Target/SPIRV/SPIRVLegalizeZeroSizeArrays.cpp
index c7f993be24636..ac51aa9ea93a8 100644
--- a/llvm/lib/Target/SPIRV/SPIRVLegalizeZeroSizeArrays.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVLegalizeZeroSizeArrays.cpp
@@ -300,10 +300,6 @@ bool SPIRVLegalizeZeroSizeArraysImpl::runOnModule(Module &M) {
if (!hasZeroSizeArray(GV.getValueType()))
continue;
- // llvm.embedded.module is handled by SPIRVPrepareGlobals.
- if (GV.getName() == "llvm.embedded.module")
- continue;
-
Type *NewTy = legalizeType(GV.getValueType());
Constant *LegalizedInitializer = legalizeConstant(GV.getInitializer());
>From 0f34421c49906ea27b30128d98f94bf5c4d745cd Mon Sep 17 00:00:00 2001
From: Alex Voicu <alexandru.voicu at amd.com>
Date: Tue, 20 Jan 2026 11:23:55 +0000
Subject: [PATCH 6/6] Update test.
---
.../CodeGen/SPIRV/fembed-bitcode-marker.ll | 21 +++++++++++--------
1 file changed, 12 insertions(+), 9 deletions(-)
diff --git a/llvm/test/CodeGen/SPIRV/fembed-bitcode-marker.ll b/llvm/test/CodeGen/SPIRV/fembed-bitcode-marker.ll
index b17464199c720..4c29a8db8c743 100644
--- a/llvm/test/CodeGen/SPIRV/fembed-bitcode-marker.ll
+++ b/llvm/test/CodeGen/SPIRV/fembed-bitcode-marker.ll
@@ -1,6 +1,6 @@
-; Expanding the bitcode marker works only for AMD at the moment.
-; RUN: not llc -verify-machineinstrs -mtriple=spirv-unknown-unknown %s -o -
-; RUN: llc -verify-machineinstrs -mtriple=spirv64-amd-amdhsa %s -o - | FileCheck %s
+; RUN: llc -verify-machineinstrs -mtriple=spirv-unknown-unknown %s -o - | FileCheck --check-prefixes=CHECK,SPIRV %s
+; RUN: llc -verify-machineinstrs -mtriple=spirv64-amd-amdhsa %s -o - | FileCheck --check-prefixes=CHECK,AMDGCNSPIRV %s
+; RUN: %if spirv-tools %{ llc -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
; RUN: %if spirv-tools %{ llc -mtriple=spirv64-amd-amdhsa %s -o - -filetype=obj | spirv-val %}
;
; Verify that we lower the embedded bitcode
@@ -11,12 +11,15 @@
; CHECK: OpName %[[#LLVM_EMBEDDED_MODULE:]] "llvm.embedded.module"
; CHECK: OpDecorate %[[#LLVM_EMBEDDED_MODULE]] Constant
; CHECK: %[[#UCHAR:]] = OpTypeInt 8 0
-; CHECK: %[[#UINT64:]] = OpTypeInt 64 0
-; CHECK: %[[#UINT64_MAX:]] = OpConstant %[[#UINT64]] 18446744073709551615
-; CHECK: %[[#UCHAR_ARR_UINT64_MAX:]] = OpTypeArray %[[#UCHAR]] %[[#UINT64_MAX]]
-; CHECK: %[[#UCHAR_ARR_UINT64_MAX_PTR:]] = OpTypePointer CrossWorkgroup %[[#UCHAR_ARR_UINT64_MAX]]
-; CHECK: %[[#CONST_UCHAR_ARR_UINT64_MAX:]] = OpConstantNull %[[#UCHAR_ARR_UINT64_MAX]]
-; CHECK: %[[#LLVM_EMBEDDED_MODULE]] = OpVariable %[[#UCHAR_ARR_UINT64_MAX_PTR]] CrossWorkgroup %[[#CONST_UCHAR_ARR_UINT64_MAX]]
+; AMDGCNSPIRV: %[[#UINT64:]] = OpTypeInt 64 0
+; SPIRV: %[[#UCHAR_PTR:]] = OpTypePointer Generic %[[#UCHAR]]
+; AMDGCNSPIRV: %[[#UINT64_MAX:]] = OpConstant %[[#UINT64]] 18446744073709551615
+; AMDGCNSPIRV: %[[#UCHAR_ARR_UINT64_MAX:]] = OpTypeArray %[[#UCHAR]] %[[#UINT64_MAX]]
+; AMDGCNSPIRV: %[[#UCHAR_ARR_UINT64_MAX_PTR:]] = OpTypePointer CrossWorkgroup %[[#UCHAR_ARR_UINT64_MAX]]
+; AMDGCNSPIRV: %[[#CONST_UCHAR_ARR_UINT64_MAX:]] = OpConstantNull %[[#UCHAR_ARR_UINT64_MAX]]
+; SPIRV: %[[#CONST_UCHAR_NULL_PTR:]] = OpConstantNull %[[#UCHAR_PTR]]
+; AMDGCNSPIRV: %[[#LLVM_EMBEDDED_MODULE]] = OpVariable %[[#UCHAR_ARR_UINT64_MAX_PTR]] CrossWorkgroup %[[#CONST_UCHAR_ARR_UINT64_MAX]]
+; SPIRV: %[[#LLVM_EMBEDDED_MODULE]] = OpVariable %[[#]] CrossWorkgroup %[[#CONST_UCHAR_NULL_PTR]]
define spir_kernel void @foo() {
entry:
More information about the llvm-commits
mailing list