[llvm] [SPIRV] Do not use OpTypeRuntimeArray in Kernel env. (PR #149522)
Marcos Maronas via llvm-commits
llvm-commits at lists.llvm.org
Tue Jul 29 06:23:23 PDT 2025
https://github.com/maarquitos14 updated https://github.com/llvm/llvm-project/pull/149522
>From 2933fcf175c25242df1c9bbbcce26aa6a73a4ae3 Mon Sep 17 00:00:00 2001
From: Marcos Maronas <marcos.maronas at intel.com>
Date: Fri, 18 Jul 2025 16:33:55 +0200
Subject: [PATCH 1/6] [SPIRV] Use OpTypeArray for Kernel env.
---
llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp | 8 ++++---
llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp | 9 +++++++-
llvm/test/CodeGen/SPIRV/array_type.ll | 22 +++++++++++++++++++
3 files changed, 35 insertions(+), 4 deletions(-)
create mode 100644 llvm/test/CodeGen/SPIRV/array_type.ll
diff --git a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp b/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
index 83fccdc2bdba3..982d48f2a5a76 100644
--- a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
@@ -828,9 +828,11 @@ SPIRVType *SPIRVGlobalRegistry::getOpTypeArray(uint32_t NumElems,
"Invalid array element type");
SPIRVType *SpvTypeInt32 = getOrCreateSPIRVIntegerType(32, MIRBuilder);
SPIRVType *ArrayType = nullptr;
- if (NumElems != 0) {
- Register NumElementsVReg =
- buildConstantInt(NumElems, MIRBuilder, SpvTypeInt32, EmitIR);
+ const SPIRVSubtarget &ST =
+ cast<SPIRVSubtarget>(MIRBuilder.getMF().getSubtarget());
+ if (NumElems != 0 || !ST.isShader()) {
+ Register NumElementsVReg = buildConstantInt(
+ NumElems ? NumElems : 1, MIRBuilder, SpvTypeInt32, EmitIR);
ArrayType = createOpType(MIRBuilder, [&](MachineIRBuilder &MIRBuilder) {
return MIRBuilder.buildInstr(SPIRV::OpTypeArray)
.addDef(createTypeVReg(MIRBuilder))
diff --git a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
index ad976e5288927..dd11b0e6bc891 100644
--- a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
@@ -68,6 +68,7 @@ getSymbolicOperandRequirements(SPIRV::OperandCategory::OperandCategory Category,
SPIRV::RequirementHandler &Reqs) {
// A set of capabilities to avoid if there is another option.
AvoidCapabilitiesSet AvoidCaps;
+
if (!ST.isShader())
AvoidCaps.S.insert(SPIRV::Capability::Shader);
else
@@ -744,8 +745,14 @@ void SPIRV::RequirementHandler::checkSatisfiable(
IsSatisfiable = false;
}
+ AvoidCapabilitiesSet AvoidCaps;
+ if (!ST.isShader())
+ AvoidCaps.S.insert(SPIRV::Capability::Shader);
+ else
+ AvoidCaps.S.insert(SPIRV::Capability::Kernel);
+
for (auto Cap : MinimalCaps) {
- if (AvailableCaps.contains(Cap))
+ if (AvailableCaps.contains(Cap) && !AvoidCaps.S.contains(Cap))
continue;
LLVM_DEBUG(dbgs() << "Capability not supported: "
<< getSymbolicOperandMnemonic(
diff --git a/llvm/test/CodeGen/SPIRV/array_type.ll b/llvm/test/CodeGen/SPIRV/array_type.ll
new file mode 100644
index 0000000000000..0957a7d191922
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/array_type.ll
@@ -0,0 +1,22 @@
+; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-KERNEL
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+
+; CHECK-KERNEL: OpCapability Kernel
+; CHECK-KERNEL-NOT: OpCapability Shader
+; CHECK-KERNEL: OpTypeArray
+; CHECK-KERNEL-NOT: OpTypeRuntimeArray
+
+%"class.sycl::_V1::detail::half_impl::half" = type { half }
+
+; Function Attrs: mustprogress norecurse nounwind
+define spir_kernel void @foo(ptr addrspace(3) noundef align 2 %_arg_temp, ptr addrspace(1) noundef align 2 %_arg_acc_a){
+entry:
+ %0 = getelementptr %"class.sycl::_V1::detail::half_impl::half", ptr addrspace(1) %_arg_acc_a, i64 15
+ %add.ptr.i = getelementptr %"class.sycl::_V1::detail::half_impl::half", ptr addrspace(1) %0, i64 10
+ %4 = getelementptr %"class.sycl::_V1::detail::half_impl::half", ptr addrspace(1) %add.ptr.i, i64 20
+ %arrayidx.i5.i = getelementptr %"class.sycl::_V1::detail::half_impl::half", ptr addrspace(1) %4, i64 35
+ %arrayidx7.i = getelementptr inbounds [0 x [32 x %"class.sycl::_V1::detail::half_impl::half"]], ptr addrspace(3) %_arg_temp, i64 1, i64 25, i64 30
+ %5 = load i16, ptr addrspace(1) %arrayidx.i5.i, align 2
+ store i16 %5, ptr addrspace(3) %arrayidx7.i, align 2
+ ret void
+}
>From 68b8e0c323c13598b021e4ef95fba20278886e86 Mon Sep 17 00:00:00 2001
From: Marcos Maronas <marcos.maronas at intel.com>
Date: Fri, 18 Jul 2025 16:38:29 +0200
Subject: [PATCH 2/6] Remove undesired change.
---
llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp | 1 -
1 file changed, 1 deletion(-)
diff --git a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
index dd11b0e6bc891..07628c6885b81 100644
--- a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp
@@ -68,7 +68,6 @@ getSymbolicOperandRequirements(SPIRV::OperandCategory::OperandCategory Category,
SPIRV::RequirementHandler &Reqs) {
// A set of capabilities to avoid if there is another option.
AvoidCapabilitiesSet AvoidCaps;
-
if (!ST.isShader())
AvoidCaps.S.insert(SPIRV::Capability::Shader);
else
>From b01c325c95f8be300838afa8a9dfe4203ebb1bd7 Mon Sep 17 00:00:00 2001
From: Marcos Maronas <marcos.maronas at intel.com>
Date: Fri, 25 Jul 2025 13:07:18 +0200
Subject: [PATCH 3/6] Simplify gep if type is 0-length array and first index is
0.
---
llvm/lib/Analysis/InstructionSimplify.cpp | 21 +++++++++++++++++++++
llvm/test/Transforms/InstSimplify/gep.ll | 17 +++++++++++++++++
2 files changed, 38 insertions(+)
diff --git a/llvm/lib/Analysis/InstructionSimplify.cpp b/llvm/lib/Analysis/InstructionSimplify.cpp
index 82530e7d5b6c6..ef77cca12cfe2 100644
--- a/llvm/lib/Analysis/InstructionSimplify.cpp
+++ b/llvm/lib/Analysis/InstructionSimplify.cpp
@@ -5013,6 +5013,27 @@ static Value *simplifyGEPInst(Type *SrcTy, Value *Ptr,
if (Indices.empty())
return Ptr;
+ // getelementptr [0 x T], P, 0 (zero), I -> getelementptr T, P, I.
+ // If type is 0-length array and first index is 0 (zero), drop both the
+ // 0-length array type and the first index. This is a common pattern in the
+ // IR, e.g. when using a zero-length array as a placeholder for a flexible
+ // array such as unbound arrays.
+ if (SrcTy->isArrayTy() && cast<ArrayType>(SrcTy)->getNumElements() == 0 &&
+ match(Indices[0], m_Zero())) {
+ assert(Q.CxtI);
+ IRBuilder<> Builder(const_cast<Instruction *>(Q.CxtI));
+ Indices = Indices.drop_front();
+ SrcTy = cast<ArrayType>(SrcTy)->getElementType();
+ Value *NewGEP = Builder.CreateGEP(SrcTy, Ptr, Indices, "", NW);
+ // Try and simplify again, it could be that after the first simplification
+ // we unlocked new simplification opportunities.
+ Value *Simplify = simplifyGEPInst(SrcTy, Ptr, Indices, NW, Q);
+ if (Simplify)
+ return Simplify;
+
+ return NewGEP;
+ }
+
// Compute the (pointer) type returned by the GEP instruction.
Type *LastType = GetElementPtrInst::getIndexedType(SrcTy, Indices);
Type *GEPTy = Ptr->getType();
diff --git a/llvm/test/Transforms/InstSimplify/gep.ll b/llvm/test/Transforms/InstSimplify/gep.ll
index 272067c66cf9f..796525622cae6 100644
--- a/llvm/test/Transforms/InstSimplify/gep.ll
+++ b/llvm/test/Transforms/InstSimplify/gep.ll
@@ -449,3 +449,20 @@ define ptr @gep_noinbounds_null(i64 %idx) {
%p = getelementptr i8, ptr null, i64 %idx
ret ptr %p
}
+
+define ptr @gep_zero_length_array_first_index_zero(i64 %idx) {
+; CHECK-LABEL: @gep_zero_length_array_first_index_zero(
+; CHECK-NEXT: [[P:%.*]] = getelementptr [32 x i32], ptr null, i64 [[IDX:%.*]]
+; CHECK-NEXT: ret ptr [[P]]
+;
+ %p = getelementptr [0 x [32 x i32]], ptr null, i64 0, i64 %idx
+ ret ptr %p
+}
+
+define ptr @gep_zero_length_array_all_index_zero(i64 %idx) {
+; CHECK-LABEL: @gep_zero_length_array_all_index_zero(
+; CHECK-NEXT: ret ptr null
+;
+ %p = getelementptr [0 x [32 x i32]], ptr null, i64 0, i64 0
+ ret ptr %p
+}
>From 1d0f71c0611e4c1e2518dc89e73c2c89c33d9658 Mon Sep 17 00:00:00 2001
From: Marcos Maronas <marcos.maronas at intel.com>
Date: Fri, 25 Jul 2025 13:10:24 +0200
Subject: [PATCH 4/6] Simplify 0-length array GEP instructions to prevent
OpTypeRuntimeArray creation in kernel mode.
---
llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp | 24 +++++-
llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp | 4 +-
llvm/test/CodeGen/SPIRV/array_type.ll | 82 ++++++++++++++++---
3 files changed, 95 insertions(+), 15 deletions(-)
diff --git a/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp b/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp
index b90e1aadbb5a1..2b80d12b09523 100644
--- a/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp
@@ -17,11 +17,13 @@
#include "SPIRVTargetMachine.h"
#include "SPIRVUtils.h"
#include "llvm/ADT/DenseSet.h"
+#include "llvm/Analysis/InstructionSimplify.h"
#include "llvm/IR/IRBuilder.h"
#include "llvm/IR/InstIterator.h"
#include "llvm/IR/InstVisitor.h"
#include "llvm/IR/IntrinsicsSPIRV.h"
#include "llvm/IR/TypedPointerType.h"
+#include "llvm/Transforms/Utils/Local.h"
#include <queue>
#include <unordered_set>
@@ -2384,14 +2386,34 @@ bool SPIRVEmitIntrinsics::runOnFunction(Function &Func) {
AggrConstTypes.clear();
AggrStores.clear();
- // fix GEP result types ahead of inference
+ // Fix GEP result types ahead of inference, and simplify if possible.
+ // Data structure for dead instructions that were simplified and replaced.
+ SmallPtrSet<Instruction *, 4> DeadInsts;
for (auto &I : instructions(Func)) {
auto *Ref = dyn_cast<GetElementPtrInst>(&I);
if (!Ref || GR->findDeducedElementType(Ref))
continue;
+ SmallVector<Value *, 8> Indices(Ref->indices());
+ Value *NewGEP = llvm::simplifyGEPInst(
+ Ref->getSourceElementType(), Ref->getPointerOperand(), Indices,
+ Ref->getNoWrapFlags(), llvm::SimplifyQuery(Func.getDataLayout(), Ref));
+ if (NewGEP) {
+ Ref->replaceAllUsesWith(NewGEP);
+ if (isInstructionTriviallyDead(Ref))
+ DeadInsts.insert(Ref);
+ if (GetElementPtrInst *NewGEPInst = dyn_cast<GetElementPtrInst>(NewGEP))
+ Ref = NewGEPInst;
+ else
+ GR->addDeducedElementType(NewGEP, normalizeType(NewGEP->getType()));
+ }
if (Type *GepTy = getGEPType(Ref))
GR->addDeducedElementType(Ref, normalizeType(GepTy));
}
+ // Remove dead instructions that were simplified and replaced.
+ for (auto *I : DeadInsts) {
+ assert(I->use_empty() && "Dead instruction should not have any uses left");
+ I->eraseFromParent();
+ }
processParamTypesByFunHeader(CurrF, B);
diff --git a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp b/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
index 982d48f2a5a76..5cbe51167f09b 100644
--- a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
@@ -830,7 +830,7 @@ SPIRVType *SPIRVGlobalRegistry::getOpTypeArray(uint32_t NumElems,
SPIRVType *ArrayType = nullptr;
const SPIRVSubtarget &ST =
cast<SPIRVSubtarget>(MIRBuilder.getMF().getSubtarget());
- if (NumElems != 0 || !ST.isShader()) {
+ if (NumElems != 0) {
Register NumElementsVReg = buildConstantInt(
NumElems ? NumElems : 1, MIRBuilder, SpvTypeInt32, EmitIR);
ArrayType = createOpType(MIRBuilder, [&](MachineIRBuilder &MIRBuilder) {
@@ -840,6 +840,8 @@ SPIRVType *SPIRVGlobalRegistry::getOpTypeArray(uint32_t NumElems,
.addUse(NumElementsVReg);
});
} else {
+ if (!ST.isShader())
+ return nullptr;
ArrayType = createOpType(MIRBuilder, [&](MachineIRBuilder &MIRBuilder) {
return MIRBuilder.buildInstr(SPIRV::OpTypeRuntimeArray)
.addDef(createTypeVReg(MIRBuilder))
diff --git a/llvm/test/CodeGen/SPIRV/array_type.ll b/llvm/test/CodeGen/SPIRV/array_type.ll
index 0957a7d191922..2db421f705a29 100644
--- a/llvm/test/CodeGen/SPIRV/array_type.ll
+++ b/llvm/test/CodeGen/SPIRV/array_type.ll
@@ -1,22 +1,78 @@
-; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-KERNEL
+; 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-KERNEL: OpCapability Kernel
-; CHECK-KERNEL-NOT: OpCapability Shader
-; CHECK-KERNEL: OpTypeArray
-; CHECK-KERNEL-NOT: OpTypeRuntimeArray
+; CHECK: OpCapability Kernel
+; CHECK-NOT: OpCapability Shader
+; CHECK-DAG: %[[#float16:]] = OpTypeFloat 16
+; CHECK-DAG: %[[#SyclHalfTy:]] = OpTypeStruct %[[#float16]]
+; CHECK-DAG: %[[#i16:]] = OpTypeInt 16
+; CHECK-DAG: %[[#i32:]] = OpTypeInt 32
+; CHECK-DAG: %[[#i64:]] = OpTypeInt 64
+; CHECK-DAG: %[[#ConstNull:]] = OpConstantNull %[[#i64]]
+; CHECK-DAG: %[[#ConstOne:]] = OpConstant %[[#i64]] 1
+; CHECK-DAG: %[[#ConstFive:]] = OpConstant %[[#i16]] 5
+; CHECK-DAG: %[[#SyclHalfTyPtr:]] = OpTypePointer Function %[[#SyclHalfTy]]
+; CHECK-DAG: %[[#i32Ptr:]] = OpTypePointer Function %[[#i32]]
+; CHECK-DAG: %[[#StorePtrTy:]] = OpTypePointer Function %[[#i16]]
%"class.sycl::_V1::detail::half_impl::half" = type { half }
; Function Attrs: mustprogress norecurse nounwind
-define spir_kernel void @foo(ptr addrspace(3) noundef align 2 %_arg_temp, ptr addrspace(1) noundef align 2 %_arg_acc_a){
+define spir_kernel void @foo(ptr %p){
+; CHECK: OpFunction
+; CHECK: %[[#Ptr:]] = OpFunctionParameter
+; CHECK: OpLabel
+; CHECK: %[[#]] = OpInBoundsPtrAccessChain %[[#SyclHalfTyPtr]] %[[#Ptr]] %[[#ConstNull]] %[[#ConstNull]]
+; CHECK: %[[#StorePtr:]] = OpBitcast %[[#StorePtrTy]] %[[#Ptr]]
+; CHECK: OpStore %[[#StorePtr]] %[[#ConstFive]]
+; CHECK: OpReturn
entry:
- %0 = getelementptr %"class.sycl::_V1::detail::half_impl::half", ptr addrspace(1) %_arg_acc_a, i64 15
- %add.ptr.i = getelementptr %"class.sycl::_V1::detail::half_impl::half", ptr addrspace(1) %0, i64 10
- %4 = getelementptr %"class.sycl::_V1::detail::half_impl::half", ptr addrspace(1) %add.ptr.i, i64 20
- %arrayidx.i5.i = getelementptr %"class.sycl::_V1::detail::half_impl::half", ptr addrspace(1) %4, i64 35
- %arrayidx7.i = getelementptr inbounds [0 x [32 x %"class.sycl::_V1::detail::half_impl::half"]], ptr addrspace(3) %_arg_temp, i64 1, i64 25, i64 30
- %5 = load i16, ptr addrspace(1) %arrayidx.i5.i, align 2
- store i16 %5, ptr addrspace(3) %arrayidx7.i, align 2
+ %0 = getelementptr inbounds [0 x [32 x %"class.sycl::_V1::detail::half_impl::half"]], ptr %p, i64 0, i64 0, i64 0
+ store i16 5, ptr %0
+ ret void
+}
+
+; Function Attrs: mustprogress norecurse nounwind
+define spir_kernel void @foo2(ptr %p){
+; CHECK: OpFunction
+; CHECK: %[[#Ptr:]] = OpFunctionParameter
+; CHECK: OpLabel
+; CHECK: %[[#BitcastOp:]] = OpInBoundsPtrAccessChain %[[#SyclHalfTyPtr]] %[[#Ptr]] %[[#ConstOne]] %[[#ConstOne]]
+; CHECK: %[[#StorePtr:]] = OpBitcast %[[#StorePtrTy]] %[[#BitcastOp]]
+; CHECK: OpStore %[[#StorePtr]] %[[#ConstFive]]
+; CHECK: OpReturn
+entry:
+ %0 = getelementptr inbounds [0 x [32 x %"class.sycl::_V1::detail::half_impl::half"]], ptr %p, i64 0, i64 1, i64 1
+ store i16 5, ptr %0
+ ret void
+}
+
+; Function Attrs: mustprogress norecurse nounwind
+define spir_kernel void @foo3(ptr %p){
+; CHECK: OpFunction
+; CHECK: %[[#Ptr:]] = OpFunctionParameter
+; CHECK: OpLabel
+; CHECK: %[[#]] = OpInBoundsPtrAccessChain %[[#i32Ptr]] %[[#Ptr]] %[[#ConstNull]] %[[#ConstNull]]
+; CHECK: %[[#StorePtr:]] = OpBitcast %[[#StorePtrTy]] %[[#Ptr]]
+; CHECK: OpStore %[[#StorePtr]] %[[#ConstFive]]
+; CHECK: OpReturn
+entry:
+ %0 = getelementptr inbounds [0 x [32 x i32]], ptr %p, i64 0, i64 0, i64 0
+ store i16 5, ptr %0
+ ret void
+}
+
+; Function Attrs: mustprogress norecurse nounwind
+define spir_kernel void @foo4(ptr %p){
+; CHECK: OpFunction
+; CHECK: %[[#Ptr:]] = OpFunctionParameter
+; CHECK: OpLabel
+; CHECK: %[[#BitcastOp:]] = OpInBoundsPtrAccessChain %[[#i32Ptr]] %[[#Ptr]] %[[#ConstOne]] %[[#ConstOne]]
+; CHECK: %[[#StorePtr:]] = OpBitcast %[[#StorePtrTy]] %[[#BitcastOp]]
+; CHECK: OpStore %[[#StorePtr]] %[[#ConstFive]]
+; CHECK: OpReturn
+entry:
+ %0 = getelementptr inbounds [0 x [32 x i32]], ptr %p, i64 0, i64 1, i64 1
+ store i16 5, ptr %0
ret void
}
>From ac4b60fca6199a582334d23ce62707bb0224b8d1 Mon Sep 17 00:00:00 2001
From: Marcos Maronas <marcos.maronas at intel.com>
Date: Fri, 25 Jul 2025 15:58:47 +0200
Subject: [PATCH 5/6] Add target environment to spir-val.
---
llvm/test/CodeGen/SPIRV/array_type.ll | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/llvm/test/CodeGen/SPIRV/array_type.ll b/llvm/test/CodeGen/SPIRV/array_type.ll
index 2db421f705a29..39d539d073462 100644
--- a/llvm/test/CodeGen/SPIRV/array_type.ll
+++ b/llvm/test/CodeGen/SPIRV/array_type.ll
@@ -1,5 +1,5 @@
; 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: %if spirv-tools %{ llc -O0 -mtriple=spirv64v1.2-unknown-unknown %s -o - -filetype=obj | spirv-val --target-env opencl2.2 %}
; CHECK: OpCapability Kernel
; CHECK-NOT: OpCapability Shader
>From b1e37f7693207d16ddf85f5b4a077dab15dd2392 Mon Sep 17 00:00:00 2001
From: Marcos Maronas <marcos.maronas at intel.com>
Date: Tue, 29 Jul 2025 15:23:03 +0200
Subject: [PATCH 6/6] Move GEP simplification from InstructionSimplify to
SPIRVEmitIntrinsics.
---
llvm/lib/Analysis/InstructionSimplify.cpp | 21 ----------
llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp | 39 ++++++++++++++-----
llvm/test/CodeGen/SPIRV/array_type.ll | 8 ++--
llvm/test/Transforms/InstSimplify/gep.ll | 17 --------
4 files changed, 34 insertions(+), 51 deletions(-)
diff --git a/llvm/lib/Analysis/InstructionSimplify.cpp b/llvm/lib/Analysis/InstructionSimplify.cpp
index ef77cca12cfe2..82530e7d5b6c6 100644
--- a/llvm/lib/Analysis/InstructionSimplify.cpp
+++ b/llvm/lib/Analysis/InstructionSimplify.cpp
@@ -5013,27 +5013,6 @@ static Value *simplifyGEPInst(Type *SrcTy, Value *Ptr,
if (Indices.empty())
return Ptr;
- // getelementptr [0 x T], P, 0 (zero), I -> getelementptr T, P, I.
- // If type is 0-length array and first index is 0 (zero), drop both the
- // 0-length array type and the first index. This is a common pattern in the
- // IR, e.g. when using a zero-length array as a placeholder for a flexible
- // array such as unbound arrays.
- if (SrcTy->isArrayTy() && cast<ArrayType>(SrcTy)->getNumElements() == 0 &&
- match(Indices[0], m_Zero())) {
- assert(Q.CxtI);
- IRBuilder<> Builder(const_cast<Instruction *>(Q.CxtI));
- Indices = Indices.drop_front();
- SrcTy = cast<ArrayType>(SrcTy)->getElementType();
- Value *NewGEP = Builder.CreateGEP(SrcTy, Ptr, Indices, "", NW);
- // Try and simplify again, it could be that after the first simplification
- // we unlocked new simplification opportunities.
- Value *Simplify = simplifyGEPInst(SrcTy, Ptr, Indices, NW, Q);
- if (Simplify)
- return Simplify;
-
- return NewGEP;
- }
-
// Compute the (pointer) type returned by the GEP instruction.
Type *LastType = GetElementPtrInst::getIndexedType(SrcTy, Indices);
Type *GEPTy = Ptr->getType();
diff --git a/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp b/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp
index 2b80d12b09523..0bdddb082986b 100644
--- a/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp
@@ -17,11 +17,11 @@
#include "SPIRVTargetMachine.h"
#include "SPIRVUtils.h"
#include "llvm/ADT/DenseSet.h"
-#include "llvm/Analysis/InstructionSimplify.h"
#include "llvm/IR/IRBuilder.h"
#include "llvm/IR/InstIterator.h"
#include "llvm/IR/InstVisitor.h"
#include "llvm/IR/IntrinsicsSPIRV.h"
+#include "llvm/IR/PatternMatch.h"
#include "llvm/IR/TypedPointerType.h"
#include "llvm/Transforms/Utils/Local.h"
@@ -189,6 +189,8 @@ class SPIRVEmitIntrinsics
void applyDemangledPtrArgTypes(IRBuilder<> &B);
+ GetElementPtrInst *simplifyZeroLengthArrayGepInst(GetElementPtrInst *GEP);
+
bool runOnFunction(Function &F);
bool postprocessTypes(Module &M);
bool processFunctionPointers(Module &M);
@@ -2369,6 +2371,29 @@ void SPIRVEmitIntrinsics::applyDemangledPtrArgTypes(IRBuilder<> &B) {
}
}
+GetElementPtrInst *
+SPIRVEmitIntrinsics::simplifyZeroLengthArrayGepInst(GetElementPtrInst *GEP) {
+ // getelementptr [0 x T], P, 0 (zero), I -> getelementptr T, P, I.
+ // If type is 0-length array and first index is 0 (zero), drop both the
+ // 0-length array type and the first index. This is a common pattern in the
+ // IR, e.g. when using a zero-length array as a placeholder for a flexible
+ // array such as unbound arrays.
+ assert(GEP && "GEP is null");
+ Type *SrcTy = GEP->getSourceElementType();
+ SmallVector<Value *, 8> Indices(GEP->indices());
+ if (SrcTy->isArrayTy() && cast<ArrayType>(SrcTy)->getNumElements() == 0 &&
+ PatternMatch::match(Indices[0], PatternMatch::m_Zero())) {
+ IRBuilder<> Builder(GEP);
+ Indices.erase(Indices.begin());
+ SrcTy = cast<ArrayType>(SrcTy)->getElementType();
+ Value *NewGEP = Builder.CreateGEP(SrcTy, GEP->getPointerOperand(), Indices,
+ "", GEP->getNoWrapFlags());
+ assert(llvm::isa<GetElementPtrInst>(NewGEP) && "NewGEP should be a GEP");
+ return cast<GetElementPtrInst>(NewGEP);
+ }
+ return nullptr;
+}
+
bool SPIRVEmitIntrinsics::runOnFunction(Function &Func) {
if (Func.isDeclaration())
return false;
@@ -2393,18 +2418,14 @@ bool SPIRVEmitIntrinsics::runOnFunction(Function &Func) {
auto *Ref = dyn_cast<GetElementPtrInst>(&I);
if (!Ref || GR->findDeducedElementType(Ref))
continue;
- SmallVector<Value *, 8> Indices(Ref->indices());
- Value *NewGEP = llvm::simplifyGEPInst(
- Ref->getSourceElementType(), Ref->getPointerOperand(), Indices,
- Ref->getNoWrapFlags(), llvm::SimplifyQuery(Func.getDataLayout(), Ref));
+
+ GetElementPtrInst *NewGEP = simplifyZeroLengthArrayGepInst(Ref);
if (NewGEP) {
Ref->replaceAllUsesWith(NewGEP);
if (isInstructionTriviallyDead(Ref))
DeadInsts.insert(Ref);
- if (GetElementPtrInst *NewGEPInst = dyn_cast<GetElementPtrInst>(NewGEP))
- Ref = NewGEPInst;
- else
- GR->addDeducedElementType(NewGEP, normalizeType(NewGEP->getType()));
+
+ Ref = NewGEP;
}
if (Type *GepTy = getGEPType(Ref))
GR->addDeducedElementType(Ref, normalizeType(GepTy));
diff --git a/llvm/test/CodeGen/SPIRV/array_type.ll b/llvm/test/CodeGen/SPIRV/array_type.ll
index 39d539d073462..251b48f8bf629 100644
--- a/llvm/test/CodeGen/SPIRV/array_type.ll
+++ b/llvm/test/CodeGen/SPIRV/array_type.ll
@@ -22,8 +22,8 @@ define spir_kernel void @foo(ptr %p){
; CHECK: OpFunction
; CHECK: %[[#Ptr:]] = OpFunctionParameter
; CHECK: OpLabel
-; CHECK: %[[#]] = OpInBoundsPtrAccessChain %[[#SyclHalfTyPtr]] %[[#Ptr]] %[[#ConstNull]] %[[#ConstNull]]
-; CHECK: %[[#StorePtr:]] = OpBitcast %[[#StorePtrTy]] %[[#Ptr]]
+; CHECK: %[[#BitcastOp:]] = OpInBoundsPtrAccessChain %[[#SyclHalfTyPtr]] %[[#Ptr]] %[[#ConstNull]] %[[#ConstNull]]
+; CHECK: %[[#StorePtr:]] = OpBitcast %[[#StorePtrTy]] %[[#BitcastOp]]
; CHECK: OpStore %[[#StorePtr]] %[[#ConstFive]]
; CHECK: OpReturn
entry:
@@ -52,8 +52,8 @@ define spir_kernel void @foo3(ptr %p){
; CHECK: OpFunction
; CHECK: %[[#Ptr:]] = OpFunctionParameter
; CHECK: OpLabel
-; CHECK: %[[#]] = OpInBoundsPtrAccessChain %[[#i32Ptr]] %[[#Ptr]] %[[#ConstNull]] %[[#ConstNull]]
-; CHECK: %[[#StorePtr:]] = OpBitcast %[[#StorePtrTy]] %[[#Ptr]]
+; CHECK: %[[#BitcastOp:]] = OpInBoundsPtrAccessChain %[[#i32Ptr]] %[[#Ptr]] %[[#ConstNull]] %[[#ConstNull]]
+; CHECK: %[[#StorePtr:]] = OpBitcast %[[#StorePtrTy]] %[[#BitcastOp]]
; CHECK: OpStore %[[#StorePtr]] %[[#ConstFive]]
; CHECK: OpReturn
entry:
diff --git a/llvm/test/Transforms/InstSimplify/gep.ll b/llvm/test/Transforms/InstSimplify/gep.ll
index 796525622cae6..272067c66cf9f 100644
--- a/llvm/test/Transforms/InstSimplify/gep.ll
+++ b/llvm/test/Transforms/InstSimplify/gep.ll
@@ -449,20 +449,3 @@ define ptr @gep_noinbounds_null(i64 %idx) {
%p = getelementptr i8, ptr null, i64 %idx
ret ptr %p
}
-
-define ptr @gep_zero_length_array_first_index_zero(i64 %idx) {
-; CHECK-LABEL: @gep_zero_length_array_first_index_zero(
-; CHECK-NEXT: [[P:%.*]] = getelementptr [32 x i32], ptr null, i64 [[IDX:%.*]]
-; CHECK-NEXT: ret ptr [[P]]
-;
- %p = getelementptr [0 x [32 x i32]], ptr null, i64 0, i64 %idx
- ret ptr %p
-}
-
-define ptr @gep_zero_length_array_all_index_zero(i64 %idx) {
-; CHECK-LABEL: @gep_zero_length_array_all_index_zero(
-; CHECK-NEXT: ret ptr null
-;
- %p = getelementptr [0 x [32 x i32]], ptr null, i64 0, i64 0
- ret ptr %p
-}
More information about the llvm-commits
mailing list