[SPIRV] Enable DCE in instruction selection and update tests #168428
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Member
|
@llvm/pr-subscribers-backend-spir-v Author: Steven Perron (s-perron) ChangesThe instruction selection pass for SPIR-V now performs dead code elimination (DCE). As a consequence of this, several tests were updated to ensure their continued
These updates improve the SPIR-V backends optimization capabilities and Patch is 111.03 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/168428.diff 54 Files Affected:
diff --git a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp b/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
index 47022b3f89a8b..fff7272f85f9e 100644
--- a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
@@ -21,6 +21,7 @@
#include "SPIRVUtils.h"
#include "llvm/ADT/APInt.h"
#include "llvm/IR/Constants.h"
+#include "llvm/IR/Function.h"
#include "llvm/IR/IntrinsicInst.h"
#include "llvm/IR/Intrinsics.h"
#include "llvm/IR/IntrinsicsSPIRV.h"
@@ -223,14 +224,44 @@ SPIRVType *SPIRVGlobalRegistry::getOpTypeVoid(MachineIRBuilder &MIRBuilder) {
}
void SPIRVGlobalRegistry::invalidateMachineInstr(MachineInstr *MI) {
+
// TODO:
+
// - review other data structure wrt. possible issues related to removal
+
// of a machine instruction during instruction selection.
+
+ // Other maps that may hold MachineInstr*:
+ // - VRegToTypeMap: Clearing would require a linear search. If we are deleting
+ // type, then no registers remaining in the code should have this type. Should
+ // be safe to leave as is.
+ // - FunctionToInstr & FunctionToInstrRev: At this point, we should not be
+ // deleting functions. No need to update.
+ // - AliasInstMDMap: Would require a linear search, and the Intel Alias
+ // instruction are not instructions instruction selection will be able to
+ // remove.
+
+ const SPIRVSubtarget &ST = MI->getMF()->getSubtarget<SPIRVSubtarget>();
+ const SPIRVInstrInfo *TII = ST.getInstrInfo();
+ assert(!TII->isAliasingInstr(*MI) &&
+ "Cannot invalidate aliasing instructions.");
+ assert(MI->getOpcode() != SPIRV::OpFunction &&
+ "Cannot invalidate OpFunction.");
+
+ if (MI->getOpcode() == SPIRV::OpFunctionCall) {
+ if (const auto *F = dyn_cast<Function>(MI->getOperand(2).getGlobal())) {
+ auto It = ForwardCalls.find(F);
+ if (It != ForwardCalls.end()) {
+ It->second.erase(MI);
+ if (It->second.empty())
+ ForwardCalls.erase(It);
+ }
+ }
+ }
+
const MachineFunction *MF = MI->getMF();
auto It = LastInsertedTypeMap.find(MF);
- if (It == LastInsertedTypeMap.end())
- return;
- if (It->second == MI)
+ if (It != LastInsertedTypeMap.end() && It->second == MI)
LastInsertedTypeMap.erase(MF);
// remove from the duplicate tracker to avoid incorrect reuse
erase(MI);
diff --git a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
index fc87288a4a212..fd473a45080eb 100644
--- a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
@@ -94,6 +94,9 @@ class SPIRVInstructionSelector : public InstructionSelector {
private:
void resetVRegsType(MachineFunction &MF);
+ // New helper function for dead instruction removal
+ void removeDeadInstruction(MachineInstr &MI) const;
+ void removeOpNamesForDeadMI(MachineInstr &MI) const;
// tblgen-erated 'select' implementation, used as the initial selector for
// the patterns that don't require complex C++.
@@ -506,22 +509,193 @@ static bool isConstReg(MachineRegisterInfo *MRI, Register OpReg) {
return false;
}
+static bool intrinsicHasSideEffects(Intrinsic::ID ID) {
+ switch (ID) {
+ // Intrinsics that do not have side effects.
+ // This is not an exhaustive list and may need to be updated.
+ case Intrinsic::spv_all:
+ case Intrinsic::spv_alloca:
+ case Intrinsic::spv_any:
+ case Intrinsic::spv_bitcast:
+ case Intrinsic::spv_const_composite:
+ case Intrinsic::spv_cross:
+ case Intrinsic::spv_degrees:
+ case Intrinsic::spv_distance:
+ case Intrinsic::spv_extractelt:
+ case Intrinsic::spv_extractv:
+ case Intrinsic::spv_faceforward:
+ case Intrinsic::spv_fdot:
+ case Intrinsic::spv_firstbitlow:
+ case Intrinsic::spv_firstbitshigh:
+ case Intrinsic::spv_firstbituhigh:
+ case Intrinsic::spv_frac:
+ case Intrinsic::spv_gep:
+ case Intrinsic::spv_global_offset:
+ case Intrinsic::spv_global_size:
+ case Intrinsic::spv_group_id:
+ case Intrinsic::spv_insertelt:
+ case Intrinsic::spv_insertv:
+ case Intrinsic::spv_isinf:
+ case Intrinsic::spv_isnan:
+ case Intrinsic::spv_lerp:
+ case Intrinsic::spv_length:
+ case Intrinsic::spv_normalize:
+ case Intrinsic::spv_num_subgroups:
+ case Intrinsic::spv_num_workgroups:
+ case Intrinsic::spv_ptrcast:
+ case Intrinsic::spv_radians:
+ case Intrinsic::spv_reflect:
+ case Intrinsic::spv_refract:
+ case Intrinsic::spv_resource_getpointer:
+ case Intrinsic::spv_resource_handlefrombinding:
+ case Intrinsic::spv_resource_handlefromimplicitbinding:
+ case Intrinsic::spv_resource_nonuniformindex:
+ case Intrinsic::spv_rsqrt:
+ case Intrinsic::spv_saturate:
+ case Intrinsic::spv_sdot:
+ case Intrinsic::spv_sign:
+ case Intrinsic::spv_smoothstep:
+ case Intrinsic::spv_step:
+ case Intrinsic::spv_subgroup_id:
+ case Intrinsic::spv_subgroup_local_invocation_id:
+ case Intrinsic::spv_subgroup_max_size:
+ case Intrinsic::spv_subgroup_size:
+ case Intrinsic::spv_thread_id:
+ case Intrinsic::spv_thread_id_in_group:
+ case Intrinsic::spv_udot:
+ case Intrinsic::spv_undef:
+ case Intrinsic::spv_value_md:
+ case Intrinsic::spv_workgroup_size:
+ return false;
+ default:
+ return true;
+ }
+}
+
+static bool isOpcodeWithNoSideEffects(unsigned Opcode) {
+ // TODO: This list should be generated by TableGen.
+ // Try to replace this with an opcode flag of some type to
+ // make sure that people are thinking about this when they add new opcodes.
+ switch (Opcode) {
+ case SPIRV::OpTypeVoid:
+ case SPIRV::OpTypeBool:
+ case SPIRV::OpTypeInt:
+ case SPIRV::OpTypeFloat:
+ case SPIRV::OpTypeVector:
+ case SPIRV::OpTypeMatrix:
+ case SPIRV::OpTypeImage:
+ case SPIRV::OpTypeSampler:
+ case SPIRV::OpTypeSampledImage:
+ case SPIRV::OpTypeArray:
+ case SPIRV::OpTypeRuntimeArray:
+ case SPIRV::OpTypeStruct:
+ case SPIRV::OpTypeOpaque:
+ case SPIRV::OpTypePointer:
+ case SPIRV::OpTypeFunction:
+ case SPIRV::OpTypeEvent:
+ case SPIRV::OpTypeDeviceEvent:
+ case SPIRV::OpTypeReserveId:
+ case SPIRV::OpTypeQueue:
+ case SPIRV::OpTypePipe:
+ case SPIRV::OpTypeForwardPointer:
+ case SPIRV::OpTypePipeStorage:
+ case SPIRV::OpTypeNamedBarrier:
+ case SPIRV::OpTypeAccelerationStructureNV:
+ case SPIRV::OpTypeCooperativeMatrixNV:
+ case SPIRV::OpTypeCooperativeMatrixKHR:
+ return true;
+ default:
+ return false;
+ }
+}
+
bool isDead(const MachineInstr &MI, const MachineRegisterInfo &MRI) {
+ // If there are no definitions, then assume there is some other
+ // side-effect that makes this instruction live.
+ if (MI.getNumDefs() == 0) {
+ return false;
+ }
+
for (const auto &MO : MI.all_defs()) {
Register Reg = MO.getReg();
- if (Reg.isPhysical() || !MRI.use_nodbg_empty(Reg))
+ if (Reg.isPhysical()) {
+ LLVM_DEBUG(dbgs() << "Not dead: def of physical register " << Reg);
return false;
+ }
+ for (const auto &UseMI : MRI.use_nodbg_instructions(Reg)) {
+ if (UseMI.getOpcode() != SPIRV::OpName) {
+ LLVM_DEBUG(dbgs() << "Not dead: def " << MO << " has use in " << UseMI);
+ return false;
+ }
+ }
}
+
if (MI.getOpcode() == TargetOpcode::LOCAL_ESCAPE || MI.isFakeUse() ||
- MI.isLifetimeMarker())
+ MI.isLifetimeMarker()) {
+ LLVM_DEBUG(
+ dbgs()
+ << "Not dead: Opcode is LOCAL_ESCAPE, fake use, or lifetime marker.\n");
return false;
- if (MI.isPHI())
+ }
+ if (MI.isPHI()) {
+ LLVM_DEBUG(dbgs() << "Dead: Phi instruction with no uses.\n");
return true;
+ }
+
+ if (MI.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
+ MI.getOpcode() == TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS) {
+ const auto &Intr = cast<GIntrinsic>(MI);
+ if (!intrinsicHasSideEffects(Intr.getIntrinsicID())) {
+ LLVM_DEBUG(dbgs() << "Dead: Intrinsic with no real side effects.\n");
+ return true;
+ }
+ }
+
if (MI.mayStore() || MI.isCall() ||
(MI.mayLoad() && MI.hasOrderedMemoryRef()) || MI.isPosition() ||
- MI.isDebugInstr() || MI.isTerminator() || MI.isJumpTableDebugInfo())
+ MI.isDebugInstr() || MI.isTerminator() || MI.isJumpTableDebugInfo()) {
+ LLVM_DEBUG(dbgs() << "Not dead: instruction has side effects.\n");
return false;
- return true;
+ }
+
+ if (isPreISelGenericOpcode(MI.getOpcode())) {
+ // TODO: Is there a generic way to check if the opcode has side effects?
+ LLVM_DEBUG(dbgs() << "Dead: Generic opcode with no uses.\n");
+ return true;
+ }
+
+ if (isOpcodeWithNoSideEffects(MI.getOpcode())) {
+ LLVM_DEBUG(dbgs() << "Dead: known opcode with no side effects\n");
+ return true;
+ }
+
+ return false;
+}
+
+void SPIRVInstructionSelector::removeOpNamesForDeadMI(MachineInstr &MI) const {
+ // Delete the OpName that uses the result of there is one.
+ for (const auto &MO : MI.all_defs()) {
+ Register Reg = MO.getReg();
+ if (Reg.isPhysical())
+ continue;
+ SmallVector<MachineInstr *, 4> UselessOpNames;
+ for (MachineInstr &UseMI : MRI->use_nodbg_instructions(Reg)) {
+ assert(UseMI.getOpcode() == SPIRV::OpName &&
+ "There is still a use of the dead function.");
+ UselessOpNames.push_back(&UseMI);
+ }
+ for (MachineInstr *OpNameMI : UselessOpNames) {
+ GR.invalidateMachineInstr(OpNameMI);
+ OpNameMI->eraseFromParent();
+ }
+ }
+}
+
+void SPIRVInstructionSelector::removeDeadInstruction(MachineInstr &MI) const {
+ salvageDebugInfo(*MRI, MI);
+ GR.invalidateMachineInstr(&MI);
+ removeOpNamesForDeadMI(MI);
+ MI.eraseFromParent();
}
bool SPIRVInstructionSelector::select(MachineInstr &I) {
@@ -530,6 +704,13 @@ bool SPIRVInstructionSelector::select(MachineInstr &I) {
assert(I.getParent() && "Instruction should be in a basic block!");
assert(I.getParent()->getParent() && "Instruction should be in a function!");
+ LLVM_DEBUG(dbgs() << "Checking if instruction is dead: " << I;);
+ if (isDead(I, *MRI)) {
+ LLVM_DEBUG(dbgs() << "Instruction is dead.\n");
+ removeDeadInstruction(I);
+ return true;
+ }
+
Register Opcode = I.getOpcode();
// If it's not a GMIR instruction, we've selected it already.
if (!isPreISelGenericOpcode(Opcode)) {
@@ -581,9 +762,7 @@ bool SPIRVInstructionSelector::select(MachineInstr &I) {
// if the instruction has been already made dead by folding it away
// erase it
LLVM_DEBUG(dbgs() << "Instruction is folded and dead.\n");
- salvageDebugInfo(*MRI, I);
- GR.invalidateMachineInstr(&I);
- I.eraseFromParent();
+ removeDeadInstruction(I);
return true;
}
diff --git a/llvm/test/CodeGen/SPIRV/OpVariable_order.ll b/llvm/test/CodeGen/SPIRV/OpVariable_order.ll
index 1e94be0886307..a43a4d66d04bb 100644
--- a/llvm/test/CodeGen/SPIRV/OpVariable_order.ll
+++ b/llvm/test/CodeGen/SPIRV/OpVariable_order.ll
@@ -13,7 +13,9 @@
define void @main() {
entry:
%0 = alloca <2 x i32>, align 4
+ store <2 x i32> zeroinitializer, ptr %0, align 4
%1 = getelementptr <2 x i32>, ptr %0, i32 0, i32 0
%2 = alloca float, align 4
+ store float 0.0, ptr %2, align 4
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/SpecConstants/restore-spec-type.ll b/llvm/test/CodeGen/SPIRV/SpecConstants/restore-spec-type.ll
index 9e91854de1172..b0bad1819a25d 100644
--- a/llvm/test/CodeGen/SPIRV/SpecConstants/restore-spec-type.ll
+++ b/llvm/test/CodeGen/SPIRV/SpecConstants/restore-spec-type.ll
@@ -29,9 +29,12 @@
%Struct7 = type [2 x %Struct]
%Nested = type { %Struct7 }
+@G = global %Struct zeroinitializer
+
define spir_kernel void @foo(ptr addrspace(4) %arg1, ptr addrspace(4) %arg2) {
entry:
%var = alloca %Struct
+ store %Struct zeroinitializer, ptr %var
%r1 = call %Struct @_Z29__spirv_SpecConstantComposite_1(float 1.0)
store %Struct %r1, ptr addrspace(4) %arg1
%r2 = call %Struct7 @_Z29__spirv_SpecConstantComposite_2(%Struct %r1, %Struct %r1)
diff --git a/llvm/test/CodeGen/SPIRV/basic_float_types.ll b/llvm/test/CodeGen/SPIRV/basic_float_types.ll
index a0ba97e1d1f14..6cdc67bbf24ee 100644
--- a/llvm/test/CodeGen/SPIRV/basic_float_types.ll
+++ b/llvm/test/CodeGen/SPIRV/basic_float_types.ll
@@ -2,6 +2,9 @@
; RUN: llc -O0 -mtriple=spirv64-unknown-unknown --spirv-ext=+SPV_KHR_bfloat16 %s -o - | FileCheck %s
; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv-unknown-unknown --spirv-ext=+SPV_KHR_bfloat16 %s -o - -filetype=obj | spirv-val %}
+// TODO: Open bug bfloat16 cannot be stored to.
+XFAIL: *
+
define void @main() {
entry:
@@ -49,50 +52,66 @@ entry:
; CHECK: %[[#]] = OpVariable %[[#ptr_Function_half]] Function
%half_Val = alloca half, align 2
+ store half 0.0, ptr %half_Val, align 2
; CHECK: %[[#]] = OpVariable %[[#ptr_Function_bfloat]] Function
%bfloat_Val = alloca bfloat, align 2
+ store bfloat 0.0, ptr %bfloat_Val, align 2
; CHECK: %[[#]] = OpVariable %[[#ptr_Function_float]] Function
%float_Val = alloca float, align 4
+ store float 0.0, ptr %float_Val, align 4
; CHECK: %[[#]] = OpVariable %[[#ptr_Function_double]] Function
%double_Val = alloca double, align 8
+ store double 0.0, ptr %double_Val, align 8
; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v2half]] Function
%half2_Val = alloca <2 x half>, align 4
+ store <2 x half> zeroinitializer, ptr %half2_Val, align 4
; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v3half]] Function
%half3_Val = alloca <3 x half>, align 8
+ store <3 x half> zeroinitializer, ptr %half3_Val, align 8
; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v4half]] Function
%half4_Val = alloca <4 x half>, align 8
+ store <4 x half> zeroinitializer, ptr %half4_Val, align 8
; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v2bfloat]] Function
%bfloat2_Val = alloca <2 x bfloat>, align 4
+ store <2 x bfloat> zeroinitializer, ptr %bfloat2_Val, align 4
; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v3bfloat]] Function
%bfloat3_Val = alloca <3 x bfloat>, align 8
+ store <3 x bfloat> zeroinitializer, ptr %bfloat3_Val, align 8
; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v4bfloat]] Function
%bfloat4_Val = alloca <4 x bfloat>, align 8
+ store <4 x bfloat> zeroinitializer, ptr %bfloat4_Val, align 8
; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v2float]] Function
%float2_Val = alloca <2 x float>, align 8
+ store <2 x float> zeroinitializer, ptr %float2_Val, align 8
; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v3float]] Function
%float3_Val = alloca <3 x float>, align 16
+ store <3 x float> zeroinitializer, ptr %float3_Val, align 16
; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v4float]] Function
%float4_Val = alloca <4 x float>, align 16
+ store <4 x float> zeroinitializer, ptr %float4_Val, align 16
; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v2double]] Function
%double2_Val = alloca <2 x double>, align 16
+ store <2 x double> zeroinitializer, ptr %double2_Val, align 16
; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v3double]] Function
%double3_Val = alloca <3 x double>, align 32
+ store <3 x double> zeroinitializer, ptr %double3_Val, align 32
; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v4double]] Function
%double4_Val = alloca <4 x double>, align 32
+ store <4 x double> zeroinitializer, ptr %double4_Val, align 32
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/basic_int_types.ll b/llvm/test/CodeGen/SPIRV/basic_int_types.ll
index 5aa7aaf6fbd01..1ed241eed4019 100644
--- a/llvm/test/CodeGen/SPIRV/basic_int_types.ll
+++ b/llvm/test/CodeGen/SPIRV/basic_int_types.ll
@@ -37,39 +37,51 @@ entry:
; CHECK: %[[#]] = OpVariable %[[#ptr_Function_short]] Function
%int16_t_Val = alloca i16, align 2
+ store i16 0, ptr %int16_t_Val, align 2
; CHECK: %[[#]] = OpVariable %[[#ptr_Function_int]] Function
%int_Val = alloca i32, align 4
+ store i32 0, ptr %int_Val, align 4
; CHECK: %[[#]] = OpVariable %[[#ptr_Function_long]] Function
%int64_t_Val = alloca i64, align 8
+ store i64 0, ptr %int64_t_Val, align 8
; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v2short]] Function
%int16_t2_Val = alloca <2 x i16>, align 4
+ store <2 x i16> zeroinitializer, ptr %int16_t2_Val, align 4
; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v3short]] Function
%int16_t3_Val = alloca <3 x i16>, align 8
+ store <3 x i16> zeroinitializer, ptr %int16_t3_Val, align 8
; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v4short]] Function
%int16_t4_Val = alloca <4 x i16>, align 8
+ store <4 x i16> zeroinitializer, ptr %int16_t4_Val, align 8
; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v2int]] Function
%int2_Val = alloca <2 x i32>, align 8
+ store <2 x i32> zeroinitializer, ptr %int2_Val, align 8
; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v3int]] Function
%int3_Val = alloca <3 x i32>, align 16
+ store <3 x i32> zeroinitializer, ptr %int3_Val, align 16
; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v4int]] Function
%int4_Val = alloca <4 x i32>, align 16
+ store <4 x i32> zeroinitializer, ptr %int4_Val, align 16
; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v2long]] Function
%int64_t2_Val = alloca <2 x i64>, align 16
+ store <2 x i64> zeroinitializer, ptr %int64_t2_Val, align 16
; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v3long]] Function
%int64_t3_Val = alloca <3 x i64>, align 32
+ store <3 x i64> zeroinitializer, ptr %int64_t3_Val, align 32
; CHECK: %[[#]] = OpVariable %[[#ptr_Function_v4long]] Function
%int64_t4_Val = alloca <4 x i64>, align 32
+ store <4 x i64> zeroinitializer, ptr %int64_t4_Val, align 32
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/basic_int_types_spirvdis.ll b/llvm/test/CodeGen/SPIRV/basic_int_types_spirvdis.ll
index 56b5f48715533..f3c8f9967211a 100644
--- a/llvm/test/CodeGen/SPIRV/basic_int_types_spirvdis.ll
+++ b/llvm/test/CodeGen/SPIRV/basic_int_types_spirvdis.ll
@@ -6,39 +6,51 @@ define void @main() {
entry:
; CHECK: %int16_t_Val = OpVariable %_ptr_Function_ushort Function
%int16_t_Val = alloca i16, align 2
+ store i16 0, i16* %int16_t_Val, align 2
; CHECK: %int_Val = OpVariable %_ptr_Function_uint Function
%int_Val = alloca i32, align 4
+ store i32 0, i32* %int_Val, align 4
; CHECK: %int64_t_Val = OpVariable %_ptr_Function_ulong Function
%int64_t_Val = alloca i64, align 8
+ store i64 0, i64* %int64_t_Val, align 8
; CHECK: %int16_t2_Val = OpVariable %_ptr_Function_v2ushort Function
%int16_t2_Val = alloca <2 x i16>, align 4
+ store <2 x i16> zeroinitializer, <2 x i16>* %int16_t2_Val, align 4
; CHECK: %int16_t3_Val = OpVariable %_ptr_Function_v3ushort Function
%int16_t3_Val = alloca <3 x i16>, align 8
+ store <3 x i16> zeroinitializer, <3 x i16>* %int16_t3_Val, align 8
; CHECK: %int16_t4_Val = OpVariable %_ptr_Function_v4ushort Function
%int16_t4_Val = alloca <4 x i16>, align 8
+ store <4 x i16> zeroinitializer, <4 x i16>* %int16_t4_Val, align 8
; CHECK: %int2_Val = OpVariable %_ptr_Function_v2uint Function
%int2_Val = alloca <2 x i32>, align 8
+ store <2 x i32> zeroinitializer, <2 x i32>* %int2_Val, align 8
; CHECK: %int3_Val = OpVariable %_ptr_Function_v3uint Function
%int3_Val = alloca <3 x i32>, align 16
+ store <3 x i32> zeroinitializer, <3 x i32>* %int3_Val, align 16
; CHECK: %int4_Val = OpVariable %_ptr_Function_v4uint Function
%int4_Val = alloca <4 x i32>, align 16
+ store <4 x i32> zeroinitializer, <4 x i32>* %int4_Val, align 16
; CHECK: %int64_t2_Val = OpVariable %_ptr_Function_v2ulong Function
%int64_t2_Val = alloca <2 x i64>, align 16
+ store <2 x i64> zeroinitializer, <2 x i64>* %int64_t2_Val, align 16
; CHECK: %int64_t3_Val = OpVariable %_ptr_Function_v3ulong Function
%int64_t3_Val = alloca <3 x i64>, align 32
+ store <3 x i64> zeroinitializer, <3 x i64>* %int64_t3_Val, align 32
; CHECK: %int64_t4_Val = OpVariable %_ptr_Function_v4ulong Function
%int64_t4_Val = alloca <4 x i64>, align 32
+ store <4 x i64> zeroinitializer, <4 x i64>* %int64_t4_Val, align 32
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/builtin_intrinsics_32.ll b/llvm/test/CodeGen/SPIRV/builtin_intrinsics_32.ll
index 39a755e736081..bca90f4ebd151 100644
--- a/llvm/test/CodeGen/SPIRV/builtin_intrinsics_32.ll
+++ b/llvm/test/CodeGen/SPIRV/builtin_intrinsics_32.ll
@@ -33,6 +33,28 @@ target triple = "spirv32-unknown-unknown"
; CHECK: [[SubgroupId]] = OpVariable [[I32PTR]] Input
; CHECK: [[SubgroupLocalInvocationId]] = OpVariable [[I32PTR]] Input
+@G_spv_num_workgroups_0 = ...
[truncated]
|
The instruction selection pass for SPIR-V now performs dead code elimination (DCE). This change removes unused instructions, leading to more optimized SPIR-V output. As a consequence of this, several tests were updated to ensure their continued correctness and to prevent previously tested code from being optimized away. Specifically: - Many tests now store computed values into global variables to ensure they are not eliminated by DCE, allowing their code generation to be verified. - The test `keep-tracked-const.ll` was removed because it no longer tested its original intent. The check statements in this test were for constants generated when expanding a G_TRUNC instruction, which is now removed by DCE instead of being expanded. - A new test, `remove-dead-type-intrinsics.ll`, was added to confirm that dead struct types are correctly removed by the compiler. These updates improve the SPIR-V backends optimization capabilities and maintain the robustness of the test suite.
🐧 Linux x64 Test Results
|
farzonl
reviewed
Nov 18, 2025
farzonl
reviewed
Nov 18, 2025
farzonl
reviewed
Nov 18, 2025
farzonl
reviewed
Nov 18, 2025
farzonl
reviewed
Nov 18, 2025
Member
|
This LGTM, issues found were all minor. |
Keenuts
reviewed
Nov 19, 2025
Contributor
Keenuts
left a comment
Keenuts
left a comment
There was a problem hiding this comment.
overall LGTM, some minor comments & nits
| // - review other data structure wrt. possible issues related to removal | ||
| // of a machine instruction during instruction selection. | ||
| // Other maps that may hold MachineInstr*: | ||
| // - VRegToTypeMap: Clearing would require a linear search. If we are deleting |
Contributor
There was a problem hiding this comment.
Expensive check this assumption?
|
|
||
| private: | ||
| void resetVRegsType(MachineFunction &MF); | ||
| // New helper function for dead instruction removal |
Contributor
There was a problem hiding this comment.
nit: remove? comment doesn't add much value no? Or maybe explain what those 2 functions are assuming/doing?
| } | ||
|
|
||
| // TODO(168736): We should make this either a flag in tabelgen | ||
| // or reduce our dependence on the global registery, so we can remove this |
Contributor
There was a problem hiding this comment.
Suggested change
| // or reduce our dependence on the global registery, so we can remove this | |
| // or reduce our dependence on the global registry, so we can remove this |
| } | ||
|
|
||
| // TODO(168736): We should make this either a flag in tabelgen | ||
| // or reduce our dependence on the global registery, so we can remove this |
Contributor
There was a problem hiding this comment.
Suggested change
| // or reduce our dependence on the global registery, so we can remove this | |
| // or reduce our dependence on the global registry, so we can remove this |
Comment on lines
+647
to
+650
| if (MI.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS || | ||
| MI.getOpcode() == TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS) { | ||
| const auto &Intr = cast<GIntrinsic>(MI); | ||
| if (!intrinsicHasSideEffects(Intr.getIntrinsicID())) { |
Contributor
There was a problem hiding this comment.
Would it be better to pass the MI into intrinsicHasSideEffect and check the opcode then the intrinsic ID inside?
We could also have the comment explaining quickly that the _SIDE_EFFECTS has no meaning for the SPIR-V backend
| } | ||
|
|
||
| void SPIRVInstructionSelector::removeOpNamesForDeadMI(MachineInstr &MI) const { | ||
| // Delete the OpName that uses the result of there is one. |
Contributor
There was a problem hiding this comment.
Suggested change
| // Delete the OpName that uses the result of there is one. | |
| // Delete the OpName that uses the result if there is one. |
farzonl
approved these changes
Nov 20, 2025
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
4 participants
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
The instruction selection pass for SPIR-V now performs dead code elimination (DCE).
This change removes unused instructions, leading to more optimized SPIR-V output.
As a consequence of this, several tests were updated to ensure their continued
correctness and to prevent previously tested code from being optimized away.
Specifically:
not eliminated by DCE, allowing their code generation to be verified.
keep-tracked-const.llwas removed because it no longer testedits original intent. The check statements in this test were for constants
generated when expanding a G_TRUNC instruction, which is now removed by DCE
instead of being expanded.
remove-dead-type-intrinsics.ll, was added to confirm that deadstruct types are correctly removed by the compiler.
These updates improve the SPIR-V backends optimization capabilities and
maintain the robustness of the test suite.