[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 07:42:21 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/9] [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/9] 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/9] 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/9] 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/9] 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/9] 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
-}
>From c724cab695607116e7dfac075150b4850158d21d Mon Sep 17 00:00:00 2001
From: Marcos Maronas <marcos.maronas at intel.com>
Date: Tue, 29 Jul 2025 16:27:42 +0200
Subject: [PATCH 7/9] Revert undesired change.
---
llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp b/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
index 5cbe51167f09b..16f6e5a19009e 100644
--- a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
@@ -832,7 +832,7 @@ SPIRVType *SPIRVGlobalRegistry::getOpTypeArray(uint32_t NumElems,
cast<SPIRVSubtarget>(MIRBuilder.getMF().getSubtarget());
if (NumElems != 0) {
Register NumElementsVReg = buildConstantInt(
- NumElems ? NumElems : 1, MIRBuilder, SpvTypeInt32, EmitIR);
+ NumElems, MIRBuilder, SpvTypeInt32, EmitIR);
ArrayType = createOpType(MIRBuilder, [&](MachineIRBuilder &MIRBuilder) {
return MIRBuilder.buildInstr(SPIRV::OpTypeArray)
.addDef(createTypeVReg(MIRBuilder))
>From 51beec0d40b625b9415e563fbeb26e24fa9716ba Mon Sep 17 00:00:00 2001
From: Marcos Maronas <marcos.maronas at intel.com>
Date: Tue, 29 Jul 2025 16:30:47 +0200
Subject: [PATCH 8/9] Fix test failures.
---
.../extensions/SPV_KHR_bit_instructions.ll | 8 --------
.../SPV_KHR_bit_instructions_no_extension.ll | 17 +++++++++++++++++
llvm/test/CodeGen/SPIRV/memory_model_md.ll | 2 +-
.../SPIRV/pointers/global-addrspacecast.ll | 8 +++++---
.../SPIRV/pointers/variables-storage-class.ll | 8 +++++---
.../SPIRV/transcoding/OpBitReverse_i32.ll | 5 +++--
.../SPIRV/transcoding/OpBitReverse_v2i16.ll | 5 +++--
7 files changed, 34 insertions(+), 19 deletions(-)
create mode 100644 llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions_no_extension.ll
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions.ll
index 40e2aff0d755a..bf71d25948a4b 100644
--- a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions.ll
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions.ll
@@ -1,18 +1,10 @@
; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv32-unknown-unknown %s --spirv-ext=+SPV_KHR_bit_instructions -o - | FileCheck %s --check-prefix=CHECK-EXTENSION
-; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-NO-EXTENSION
; CHECK-EXTENSION: OpCapability BitInstructions
; CHECK-EXTENSION-NEXT: OpExtension "SPV_KHR_bit_instructions"
; CHECK-EXTENSION-NOT: OpCabilitity Shader
-; CHECK-NO-EXTENSION: OpCapability Shader
-; CHECK-NO-EXTENSION-NOT: OpCabilitity BitInstructions
-; CHECK-NO-EXTENSION-NOT: OpExtension "SPV_KHR_bit_instructions"
-
-
; CHECK-EXTENSION: %[[#int:]] = OpTypeInt 32
; CHECK-EXTENSION: OpBitReverse %[[#int]]
-; CHECK-NO-EXTENSION: %[[#int:]] = OpTypeInt 32
-; CHECK-NO-EXTENSION: OpBitReverse %[[#int]]
define spir_kernel void @testBitRev(i32 %a, i32 %b, i32 %c, i32 addrspace(1)* nocapture %res) local_unnamed_addr {
entry:
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions_no_extension.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions_no_extension.ll
new file mode 100644
index 0000000000000..35b7df773bf07
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bit_instructions_no_extension.ll
@@ -0,0 +1,17 @@
+; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv32-unknown-vulkan %s -o - | FileCheck %s --check-prefix=CHECK-NO-EXTENSION
+
+; CHECK-NO-EXTENSION: OpCapability Shader
+; CHECK-NO-EXTENSION-NOT: OpCabilitity BitInstructions
+; CHECK-NO-EXTENSION-NOT: OpExtension "SPV_KHR_bit_instructions"
+; CHECK-NO-EXTENSION: %[[#int:]] = OpTypeInt 32
+; CHECK-NO-EXTENSION: OpBitReverse %[[#int]]
+
+define spir_kernel void @testBitRev(i32 %a, i32 %b, i32 %c, i32 addrspace(1)* nocapture %res) local_unnamed_addr #0 {
+entry:
+ %call = tail call i32 @llvm.bitreverse.i32(i32 %b)
+ store i32 %call, i32 addrspace(1)* %res, align 4
+ ret void
+}
+
+declare i32 @llvm.bitreverse.i32(i32)
+attributes #0 = { "hlsl.shader"="compute" }
diff --git a/llvm/test/CodeGen/SPIRV/memory_model_md.ll b/llvm/test/CodeGen/SPIRV/memory_model_md.ll
index e52343cbbb7e4..b2a56e34296f1 100644
--- a/llvm/test/CodeGen/SPIRV/memory_model_md.ll
+++ b/llvm/test/CodeGen/SPIRV/memory_model_md.ll
@@ -1,4 +1,4 @@
-; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=SPV
+; RUN: llc -O0 -mtriple=spirv32-unknown-vulkan %s -o - | FileCheck %s --check-prefix=SPV
; SPV: OpMemoryModel Physical32 Simple
define dso_local dllexport void @k_no_fc(i32 %ibuf, i32 %obuf) local_unnamed_addr {
diff --git a/llvm/test/CodeGen/SPIRV/pointers/global-addrspacecast.ll b/llvm/test/CodeGen/SPIRV/pointers/global-addrspacecast.ll
index 544c657da8488..44b44fc614819 100644
--- a/llvm/test/CodeGen/SPIRV/pointers/global-addrspacecast.ll
+++ b/llvm/test/CodeGen/SPIRV/pointers/global-addrspacecast.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: llc -verify-machineinstrs -O0 -mtriple=spirv64-unknown-vulkan %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-vulkan %s -o - -filetype=obj | spirv-val %}
@PrivInternal = internal addrspace(10) global i32 456
; CHECK-DAG: %[[#type:]] = OpTypeInt 32 0
@@ -7,7 +7,7 @@
; CHECK-DAG: %[[#value:]] = OpConstant %[[#type]] 456
; CHECK-DAG: %[[#var:]] = OpVariable %[[#ptrty]] Private %[[#value]]
-define spir_kernel void @Foo() {
+define spir_kernel void @Foo() #0 {
%p = addrspacecast ptr addrspace(10) @PrivInternal to ptr
%v = load i32, ptr %p, align 4
ret void
@@ -15,3 +15,5 @@ define spir_kernel void @Foo() {
; CHECK-NEXT: OpLoad %[[#type]] %[[#var]] Aligned 4
; CHECK-Next: OpReturn
}
+
+attributes #0 = { "hlsl.shader"="compute" }
diff --git a/llvm/test/CodeGen/SPIRV/pointers/variables-storage-class.ll b/llvm/test/CodeGen/SPIRV/pointers/variables-storage-class.ll
index a1ded0569d67e..84c3f257c6b72 100644
--- a/llvm/test/CodeGen/SPIRV/pointers/variables-storage-class.ll
+++ b/llvm/test/CodeGen/SPIRV/pointers/variables-storage-class.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: llc -verify-machineinstrs -O0 -mtriple=spirv64-unknown-vulkan %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-vulkan %s -o - -filetype=obj | spirv-val %}
; CHECK-DAG: %[[#U8:]] = OpTypeInt 8 0
; CHECK-DAG: %[[#U32:]] = OpTypeInt 32 0
@@ -20,7 +20,7 @@
; CHECK-DAG: %[[#]] = OpVariable %[[#VTYPE]] Private %[[#VAL]]
@PrivInternal = internal addrspace(10) global i32 456
-define spir_kernel void @Foo() {
+define spir_kernel void @Foo() #0 {
; CHECK: %[[#]] = OpLoad %[[#]] %[[#PTR]] Aligned 8
%l = load ptr addrspace(1), ptr addrspace(1) @Ptr, align 8
; CHECK: OpCopyMemorySized %[[#]] %[[#INIT]] %[[#]] Aligned 4
@@ -29,3 +29,5 @@ define spir_kernel void @Foo() {
}
declare void @llvm.memcpy.p1.p2.i64(ptr addrspace(1) noalias nocapture writeonly, ptr addrspace(2) noalias nocapture readonly, i64, i1 immarg)
+
+attributes #0 = { "hlsl.shader"="compute" }
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_i32.ll b/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_i32.ll
index f396b5a01ae91..f8ffa98bdce81 100644
--- a/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_i32.ll
+++ b/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_i32.ll
@@ -1,9 +1,9 @@
-; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: llc -O0 -mtriple=spirv32-unknown-vulkan %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
; CHECK-SPIRV: %[[#int:]] = OpTypeInt 32
; CHECK-SPIRV: OpBitReverse %[[#int]]
-define spir_kernel void @testBitRev(i32 %a, i32 %b, i32 %c, i32 addrspace(1)* nocapture %res) local_unnamed_addr {
+define spir_kernel void @testBitRev(i32 %a, i32 %b, i32 %c, i32 addrspace(1)* nocapture %res) local_unnamed_addr #0 {
entry:
%call = tail call i32 @llvm.bitreverse.i32(i32 %b)
store i32 %call, i32 addrspace(1)* %res, align 4
@@ -11,3 +11,4 @@ entry:
}
declare i32 @llvm.bitreverse.i32(i32)
+attributes #0 = { "hlsl.shader"="compute" }
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_v2i16.ll b/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_v2i16.ll
index 8f04929fdd587..02bd41eac7568 100644
--- a/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_v2i16.ll
+++ b/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse_v2i16.ll
@@ -1,10 +1,10 @@
-; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: llc -O0 -mtriple=spirv32-unknown-vulkan %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
; CHECK-SPIRV: %[[#short:]] = OpTypeInt 16
; CHECK-SPIRV: %[[#short2:]] = OpTypeVector %[[#short]] 2
; CHECK-SPIRV: OpBitReverse %[[#short2]]
-define spir_kernel void @testBitRev(<2 x i16> %a, <2 x i16> %b, <2 x i16> %c, <2 x i16> addrspace(1)* nocapture %res) local_unnamed_addr {
+define spir_kernel void @testBitRev(<2 x i16> %a, <2 x i16> %b, <2 x i16> %c, <2 x i16> addrspace(1)* nocapture %res) local_unnamed_addr #0 {
entry:
%call = tail call <2 x i16> @llvm.bitreverse.v2i16(<2 x i16> %b)
store <2 x i16> %call, <2 x i16> addrspace(1)* %res, align 4
@@ -12,3 +12,4 @@ entry:
}
declare <2 x i16> @llvm.bitreverse.v2i16(<2 x i16>)
+attributes #0 = { "hlsl.shader"="compute" }
>From 0f4ee2b3b5cb7c96cc92c3b96a04aae123e1c76b Mon Sep 17 00:00:00 2001
From: Marcos Maronas <marcos.maronas at intel.com>
Date: Tue, 29 Jul 2025 16:40:11 +0200
Subject: [PATCH 9/9] Fix clang-format issue.
---
llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp b/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
index 16f6e5a19009e..960eb2ef93a9e 100644
--- a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
@@ -831,8 +831,8 @@ SPIRVType *SPIRVGlobalRegistry::getOpTypeArray(uint32_t NumElems,
const SPIRVSubtarget &ST =
cast<SPIRVSubtarget>(MIRBuilder.getMF().getSubtarget());
if (NumElems != 0) {
- Register NumElementsVReg = buildConstantInt(
- NumElems, MIRBuilder, SpvTypeInt32, EmitIR);
+ Register NumElementsVReg =
+ buildConstantInt(NumElems, MIRBuilder, SpvTypeInt32, EmitIR);
ArrayType = createOpType(MIRBuilder, [&](MachineIRBuilder &MIRBuilder) {
return MIRBuilder.buildInstr(SPIRV::OpTypeArray)
.addDef(createTypeVReg(MIRBuilder))
More information about the llvm-commits
mailing list