[llvm] [SPIRV] Enable DCE in instruction selection and update tests (PR #168428)
Steven Perron via llvm-commits
llvm-commits at lists.llvm.org
Thu Nov 20 09:26:18 PST 2025
https://github.com/s-perron updated https://github.com/llvm/llvm-project/pull/168428
>From 4d5bfba8bf2deeea50465b22caaedcd5f154bdc6 Mon Sep 17 00:00:00 2001
From: Steven Perron <stevenperron at google.com>
Date: Tue, 7 Oct 2025 13:26:47 -0400
Subject: [PATCH 1/4] [SPIRV] Enable DCE in instruction selection and update
tests
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.
---
llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp | 36 +++-
.../Target/SPIRV/SPIRVInstructionSelector.cpp | 195 +++++++++++++++++-
llvm/test/CodeGen/SPIRV/OpVariable_order.ll | 2 +
.../SPIRV/SpecConstants/restore-spec-type.ll | 3 +
llvm/test/CodeGen/SPIRV/basic_float_types.ll | 19 ++
llvm/test/CodeGen/SPIRV/basic_int_types.ll | 12 ++
.../CodeGen/SPIRV/basic_int_types_spirvdis.ll | 12 ++
.../CodeGen/SPIRV/builtin_intrinsics_32.ll | 43 ++++
.../CodeGen/SPIRV/builtin_intrinsics_64.ll | 43 ++++
.../CodeGen/SPIRV/builtin_vars-decorate.ll | 19 ++
.../SPIRV/debug-info/debug-type-pointer.ll | 2 +
llvm/test/CodeGen/SPIRV/event-zero-const.ll | 4 +
.../fun-ptr-addrcast.ll | 7 +
.../extensions/SPV_KHR_bfloat16/bfloat16.ll | 5 +
.../SPV_KHR_float_controls2/decoration.ll | 94 +++++++++
.../enable-all-extensions-but-one.ll | 5 +
llvm/test/CodeGen/SPIRV/freeze.ll | 48 +++--
.../SPIRV/hlsl-intrinsics/AddUint64.ll | 2 +-
.../test/CodeGen/SPIRV/hlsl-intrinsics/abs.ll | 17 +-
.../CodeGen/SPIRV/hlsl-intrinsics/log10.ll | 8 +-
.../instructions/insertvalue-undef-ptr.ll | 17 +-
.../SPIRV/instructions/select-ptr-load.ll | 5 +
llvm/test/CodeGen/SPIRV/keep-tracked-const.ll | 23 ---
.../CodeGen/SPIRV/llvm-intrinsics/assume.ll | 13 +-
.../llvm-intrinsics/bitreverse_small_type.ll | 53 +++--
.../llvm-intrinsics/constrained-arithmetic.ll | 13 ++
.../CodeGen/SPIRV/llvm-intrinsics/lifetime.ll | 22 +-
.../SPIRV/llvm-intrinsics/satur-arith.ll | 38 +++-
.../llvm-intrinsics/uadd.with.overflow.ll | 5 +-
.../CodeGen/SPIRV/logical-access-chain.ll | 5 +-
.../CodeGen/SPIRV/logical-struct-access.ll | 83 ++++++--
llvm/test/CodeGen/SPIRV/phi-insert-point.ll | 13 ++
.../CodeGen/SPIRV/phi-ptrcast-dominate.ll | 9 +
.../SPIRV/pointers/bitcast-fix-accesschain.ll | 5 +
.../SPIRV/pointers/bitcast-fix-load.ll | 3 +
.../CodeGen/SPIRV/pointers/gep-types-1.ll | 3 +
.../pointers/getelementptr-addressspace.ll | 5 +
.../SPIRV/pointers/getelementptr-base-type.ll | 3 +
.../pointers/getelementptr-bitcast-load.ll | 5 +
.../pointers/getelementptr-kernel-arg-char.ll | 5 +
.../SPIRV/pointers/global-addrspacecast.ll | 7 +-
.../SPIRV/pointers/load-addressspace.ll | 5 +
.../CodeGen/SPIRV/pointers/phi-chain-types.ll | 2 +
.../SPIRV/pointers/pointer-addrspacecast.ll | 3 +
.../CodeGen/SPIRV/pointers/ptr-eq-types.ll | 6 +
.../pointers/resource-vector-load-store.ll | 27 ++-
.../pointers/type-deduce-call-no-bitcast.ll | 3 +
.../SPIRV/remove-dead-type-intrinsics.ll | 31 +++
.../SPIRV/transcoding/OpBitReverse-subbyte.ll | 5 +
.../SPIRV/transcoding/OpGenericCastToPtr.ll | 18 ++
.../SPIRV/transcoding/OpPtrCastToGeneric.ll | 5 +
llvm/test/CodeGen/SPIRV/transcoding/fcmp.ll | 182 ++++++++++++++++
.../SPIRV/transcoding/spirv-event-null.ll | 4 +-
llvm/test/CodeGen/SPIRV/uitofp-with-bool.ll | 41 ++++
54 files changed, 1110 insertions(+), 133 deletions(-)
delete mode 100644 llvm/test/CodeGen/SPIRV/keep-tracked-const.ll
create mode 100644 llvm/test/CodeGen/SPIRV/remove-dead-type-intrinsics.ll
diff --git a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp b/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
index 47022b3f89a8b..36facbd05aa67 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,37 @@ 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 }
+ at 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
+ at G_spv_num_workgroups_0 = global i32 0
+ at G_spv_num_workgroups_1 = global i32 0
+ at G_spv_num_workgroups_2 = global i32 0
+ at G_spv_workgroup_size_0 = global i32 0
+ at G_spv_workgroup_size_1 = global i32 0
+ at G_spv_workgroup_size_2 = global i32 0
+ at G_spv_group_id_0 = global i32 0
+ at G_spv_group_id_1 = global i32 0
+ at G_spv_group_id_2 = global i32 0
+ at G_spv_thread_id_in_group_0 = global i32 0
+ at G_spv_thread_id_in_group_1 = global i32 0
+ at G_spv_thread_id_in_group_2 = global i32 0
+ at G_spv_thread_id_0 = global i32 0
+ at G_spv_thread_id_1 = global i32 0
+ at G_spv_thread_id_2 = global i32 0
+ at G_spv_global_size_0 = global i32 0
+ at G_spv_global_size_1 = global i32 0
+ at G_spv_global_size_2 = global i32 0
+ at G_spv_global_offset_0 = global i32 0
+ at G_spv_global_offset_1 = global i32 0
+ at G_spv_global_offset_2 = global i32 0
+
; Function Attrs: convergent noinline norecurse nounwind optnone
define spir_func void @test_id_and_range() {
entry:
@@ -44,66 +66,87 @@ entry:
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[NumWorkgroups]]
; CHECK: OpCompositeExtract [[I32]] [[LD]] 0
%spv.num.workgroups = call i32 @llvm.spv.num.workgroups.i32(i32 0)
+ store i32 %spv.num.workgroups, i32* @G_spv_num_workgroups_0
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[NumWorkgroups]]
; CHECK: OpCompositeExtract [[I32]] [[LD]] 1
%spv.num.workgroups1 = call i32 @llvm.spv.num.workgroups.i32(i32 1)
+ store i32 %spv.num.workgroups1, i32* @G_spv_num_workgroups_1
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[NumWorkgroups]]
; CHECK: OpCompositeExtract [[I32]] [[LD]] 2
%spv.num.workgroups2 = call i32 @llvm.spv.num.workgroups.i32(i32 2)
+ store i32 %spv.num.workgroups2, i32* @G_spv_num_workgroups_2
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[WorkgroupSize]]
; CHECK: OpCompositeExtract [[I32]] [[LD]] 0
%spv.workgroup.size = call i32 @llvm.spv.workgroup.size.i32(i32 0)
+ store i32 %spv.workgroup.size, i32* @G_spv_workgroup_size_0
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[WorkgroupSize]]
; CHECK: OpCompositeExtract [[I32]] [[LD]] 1
%spv.workgroup.size3 = call i32 @llvm.spv.workgroup.size.i32(i32 1)
+ store i32 %spv.workgroup.size3, i32* @G_spv_workgroup_size_1
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[WorkgroupSize]]
; CHECK: OpCompositeExtract [[I32]] [[LD]] 2
%spv.workgroup.size4 = call i32 @llvm.spv.workgroup.size.i32(i32 2)
+ store i32 %spv.workgroup.size4, i32* @G_spv_workgroup_size_2
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[WorkgroupId]]
; CHECK: OpCompositeExtract [[I32]] [[LD]] 0
%spv.group.id = call i32 @llvm.spv.group.id.i32(i32 0)
+ store i32 %spv.group.id, i32* @G_spv_group_id_0
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[WorkgroupId]]
; CHECK: OpCompositeExtract [[I32]] [[LD]] 1
%spv.group.id5 = call i32 @llvm.spv.group.id.i32(i32 1)
+ store i32 %spv.group.id5, i32* @G_spv_group_id_1
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[WorkgroupId]]
; CHECK: OpCompositeExtract [[I32]] [[LD]] 2
%spv.group.id6 = call i32 @llvm.spv.group.id.i32(i32 2)
+ store i32 %spv.group.id6, i32* @G_spv_group_id_2
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[LocalInvocationId]]
; CHECK: OpCompositeExtract [[I32]] [[LD]] 0
%spv.thread.id.in.group = call i32 @llvm.spv.thread.id.in.group.i32(i32 0)
+ store i32 %spv.thread.id.in.group, i32* @G_spv_thread_id_in_group_0
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[LocalInvocationId]]
; CHECK: OpCompositeExtract [[I32]] [[LD]] 1
%spv.thread.id.in.group7 = call i32 @llvm.spv.thread.id.in.group.i32(i32 1)
+ store i32 %spv.thread.id.in.group7, i32* @G_spv_thread_id_in_group_1
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[LocalInvocationId]]
; CHECK: OpCompositeExtract [[I32]] [[LD]] 2
%spv.thread.id.in.group8 = call i32 @llvm.spv.thread.id.in.group.i32(i32 2)
+ store i32 %spv.thread.id.in.group8, i32* @G_spv_thread_id_in_group_2
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[GlobalInvocationId]]
; CHECK: OpCompositeExtract [[I32]] [[LD]] 0
%spv.thread.id = call i32 @llvm.spv.thread.id.i32(i32 0)
+ store i32 %spv.thread.id, i32* @G_spv_thread_id_0
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[GlobalInvocationId]]
; CHECK: OpCompositeExtract [[I32]] [[LD]] 1
%spv.thread.id9 = call i32 @llvm.spv.thread.id.i32(i32 1)
+ store i32 %spv.thread.id9, i32* @G_spv_thread_id_1
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[GlobalInvocationId]]
; CHECK: OpCompositeExtract [[I32]] [[LD]] 2
%spv.thread.id10 = call i32 @llvm.spv.thread.id.i32(i32 2)
+ store i32 %spv.thread.id10, i32* @G_spv_thread_id_2
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[GlobalSize]]
; CHECK: OpCompositeExtract [[I32]] [[LD]] 0
%spv.num.workgroups11 = call i32 @llvm.spv.global.size.i32(i32 0)
+ store i32 %spv.num.workgroups11, i32* @G_spv_global_size_0
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[GlobalSize]]
; CHECK: OpCompositeExtract [[I32]] [[LD]] 1
%spv.num.workgroups12 = call i32 @llvm.spv.global.size.i32(i32 1)
+ store i32 %spv.num.workgroups12, i32* @G_spv_global_size_1
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[GlobalSize]]
; CHECK: OpCompositeExtract [[I32]] [[LD]] 2
%spv.num.workgroups13 = call i32 @llvm.spv.global.size.i32(i32 2)
+ store i32 %spv.num.workgroups13, i32* @G_spv_global_size_2
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[GlobalOffset]]
; CHECK: OpCompositeExtract [[I32]] [[LD]] 0
%spv.global.offset = call i32 @llvm.spv.global.offset.i32(i32 0)
+ store i32 %spv.global.offset, i32* @G_spv_global_offset_0
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[GlobalOffset]]
; CHECK: OpCompositeExtract [[I32]] [[LD]] 1
%spv.global.offset14 = call i32 @llvm.spv.global.offset.i32(i32 1)
+ store i32 %spv.global.offset14, i32* @G_spv_global_offset_1
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I32V3]] [[GlobalOffset]]
; CHECK: OpCompositeExtract [[I32]] [[LD]] 2
%spv.global.offset15 = call i32 @llvm.spv.global.offset.i32(i32 2)
+ store i32 %spv.global.offset15, i32* @G_spv_global_offset_2
; CHECK: OpLoad %5 [[SubgroupSize]]
%0 = call i32 @llvm.spv.subgroup.size()
store i32 %0, ptr %ssize, align 4
diff --git a/llvm/test/CodeGen/SPIRV/builtin_intrinsics_64.ll b/llvm/test/CodeGen/SPIRV/builtin_intrinsics_64.ll
index dcdf8992ce1c4..26c2d866d14c7 100644
--- a/llvm/test/CodeGen/SPIRV/builtin_intrinsics_64.ll
+++ b/llvm/test/CodeGen/SPIRV/builtin_intrinsics_64.ll
@@ -34,6 +34,28 @@ target triple = "spirv64-unknown-unknown"
; CHECK: [[SubgroupId]] = OpVariable [[I32PTR]] Input
; CHECK: [[SubgroupLocalInvocationId]] = OpVariable [[I32PTR]] Input
+ at G_spv_num_workgroups_0 = global i64 0
+ at G_spv_num_workgroups_1 = global i64 0
+ at G_spv_num_workgroups_2 = global i64 0
+ at G_spv_workgroup_size_0 = global i64 0
+ at G_spv_workgroup_size_1 = global i64 0
+ at G_spv_workgroup_size_2 = global i64 0
+ at G_spv_group_id_0 = global i64 0
+ at G_spv_group_id_1 = global i64 0
+ at G_spv_group_id_2 = global i64 0
+ at G_spv_thread_id_in_group_0 = global i64 0
+ at G_spv_thread_id_in_group_1 = global i64 0
+ at G_spv_thread_id_in_group_2 = global i64 0
+ at G_spv_thread_id_0 = global i64 0
+ at G_spv_thread_id_1 = global i64 0
+ at G_spv_thread_id_2 = global i64 0
+ at G_spv_global_size_0 = global i64 0
+ at G_spv_global_size_1 = global i64 0
+ at G_spv_global_size_2 = global i64 0
+ at G_spv_global_offset_0 = global i64 0
+ at G_spv_global_offset_1 = global i64 0
+ at G_spv_global_offset_2 = global i64 0
+
; Function Attrs: convergent noinline norecurse nounwind optnone
define spir_func void @test_id_and_range() {
entry:
@@ -45,66 +67,87 @@ entry:
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[NumWorkgroups]]
; CHECK: OpCompositeExtract [[I64]] [[LD]] 0
%spv.num.workgroups = call i64 @llvm.spv.num.workgroups.i64(i32 0)
+ store i64 %spv.num.workgroups, i64* @G_spv_num_workgroups_0
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[NumWorkgroups]]
; CHECK: OpCompositeExtract [[I64]] [[LD]] 1
%spv.num.workgroups1 = call i64 @llvm.spv.num.workgroups.i64(i32 1)
+ store i64 %spv.num.workgroups1, i64* @G_spv_num_workgroups_1
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[NumWorkgroups]]
; CHECK: OpCompositeExtract [[I64]] [[LD]] 2
%spv.num.workgroups2 = call i64 @llvm.spv.num.workgroups.i64(i32 2)
+ store i64 %spv.num.workgroups2, i64* @G_spv_num_workgroups_2
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[WorkgroupSize]]
; CHECK: OpCompositeExtract [[I64]] [[LD]] 0
%spv.workgroup.size = call i64 @llvm.spv.workgroup.size.i64(i32 0)
+ store i64 %spv.workgroup.size, i64* @G_spv_workgroup_size_0
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[WorkgroupSize]]
; CHECK: OpCompositeExtract [[I64]] [[LD]] 1
%spv.workgroup.size3 = call i64 @llvm.spv.workgroup.size.i64(i32 1)
+ store i64 %spv.workgroup.size3, i64* @G_spv_workgroup_size_1
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[WorkgroupSize]]
; CHECK: OpCompositeExtract [[I64]] [[LD]] 2
%spv.workgroup.size4 = call i64 @llvm.spv.workgroup.size.i64(i32 2)
+ store i64 %spv.workgroup.size4, i64* @G_spv_workgroup_size_2
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[WorkgroupId]]
; CHECK: OpCompositeExtract [[I64]] [[LD]] 0
%spv.group.id = call i64 @llvm.spv.group.id.i64(i32 0)
+ store i64 %spv.group.id, i64* @G_spv_group_id_0
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[WorkgroupId]]
; CHECK: OpCompositeExtract [[I64]] [[LD]] 1
%spv.group.id5 = call i64 @llvm.spv.group.id.i64(i32 1)
+ store i64 %spv.group.id5, i64* @G_spv_group_id_1
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[WorkgroupId]]
; CHECK: OpCompositeExtract [[I64]] [[LD]] 2
%spv.group.id6 = call i64 @llvm.spv.group.id.i64(i32 2)
+ store i64 %spv.group.id6, i64* @G_spv_group_id_2
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[LocalInvocationId]]
; CHECK: OpCompositeExtract [[I64]] [[LD]] 0
%spv.thread.id.in.group = call i64 @llvm.spv.thread.id.in.group.i64(i32 0)
+ store i64 %spv.thread.id.in.group, i64* @G_spv_thread_id_in_group_0
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[LocalInvocationId]]
; CHECK: OpCompositeExtract [[I64]] [[LD]] 1
%spv.thread.id.in.group7 = call i64 @llvm.spv.thread.id.in.group.i64(i32 1)
+ store i64 %spv.thread.id.in.group7, i64* @G_spv_thread_id_in_group_1
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[LocalInvocationId]]
; CHECK: OpCompositeExtract [[I64]] [[LD]] 2
%spv.thread.id.in.group8 = call i64 @llvm.spv.thread.id.in.group.i64(i32 2)
+ store i64 %spv.thread.id.in.group8, i64* @G_spv_thread_id_in_group_2
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[GlobalInvocationId]]
; CHECK: OpCompositeExtract [[I64]] [[LD]] 0
%spv.thread.id = call i64 @llvm.spv.thread.id.i64(i32 0)
+ store i64 %spv.thread.id, i64* @G_spv_thread_id_0
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[GlobalInvocationId]]
; CHECK: OpCompositeExtract [[I64]] [[LD]] 1
%spv.thread.id9 = call i64 @llvm.spv.thread.id.i64(i32 1)
+ store i64 %spv.thread.id9, i64* @G_spv_thread_id_1
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[GlobalInvocationId]]
; CHECK: OpCompositeExtract [[I64]] [[LD]] 2
%spv.thread.id10 = call i64 @llvm.spv.thread.id.i64(i32 2)
+ store i64 %spv.thread.id10, i64* @G_spv_thread_id_2
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[GlobalSize]]
; CHECK: OpCompositeExtract [[I64]] [[LD]] 0
%spv.num.workgroups11 = call i64 @llvm.spv.global.size.i64(i32 0)
+ store i64 %spv.num.workgroups11, i64* @G_spv_global_size_0
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[GlobalSize]]
; CHECK: OpCompositeExtract [[I64]] [[LD]] 1
%spv.num.workgroups12 = call i64 @llvm.spv.global.size.i64(i32 1)
+ store i64 %spv.num.workgroups12, i64* @G_spv_global_size_1
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[GlobalSize]]
; CHECK: OpCompositeExtract [[I64]] [[LD]] 2
%spv.num.workgroups13 = call i64 @llvm.spv.global.size.i64(i32 2)
+ store i64 %spv.num.workgroups13, i64* @G_spv_global_size_2
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[GlobalOffset]]
; CHECK: OpCompositeExtract [[I64]] [[LD]] 0
%spv.global.offset = call i64 @llvm.spv.global.offset.i64(i32 0)
+ store i64 %spv.global.offset, i64* @G_spv_global_offset_0
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[GlobalOffset]]
; CHECK: OpCompositeExtract [[I64]] [[LD]] 1
%spv.global.offset14 = call i64 @llvm.spv.global.offset.i64(i32 1)
+ store i64 %spv.global.offset14, i64* @G_spv_global_offset_1
; CHECK: [[LD:%[0-9]*]] = OpLoad [[I64V3]] [[GlobalOffset]]
; CHECK: OpCompositeExtract [[I64]] [[LD]] 2
%spv.global.offset15 = call i64 @llvm.spv.global.offset.i64(i32 2)
+ store i64 %spv.global.offset15, i64* @G_spv_global_offset_2
; CHECK: OpLoad %5 [[SubgroupSize]]
%0 = call i32 @llvm.spv.subgroup.size()
store i32 %0, ptr %ssize, align 4
diff --git a/llvm/test/CodeGen/SPIRV/builtin_vars-decorate.ll b/llvm/test/CodeGen/SPIRV/builtin_vars-decorate.ll
index 0c9b29de890d4..8dd9b387a6d84 100644
--- a/llvm/test/CodeGen/SPIRV/builtin_vars-decorate.ll
+++ b/llvm/test/CodeGen/SPIRV/builtin_vars-decorate.ll
@@ -81,17 +81,36 @@
@__spirv_BuiltInSubgroupId = external addrspace(1) global i32
@__spirv_BuiltInSubgroupLocalInvocationId = external addrspace(1) global i32
+ at G_r1 = global i64 0
+ at G_r2 = global i64 0
+ at G_r3 = global i32 0
+ at G_r4 = global i32 0
+ at G_r5 = global i32 0
+ at G_r6 = global i32 0
+ at G_r7 = global i32 0
+ at G_r8 = global i32 0
+ at G_r9 = global i32 0
+
define spir_kernel void @_Z1wv() {
entry:
%r1 = tail call spir_func i64 @get_global_linear_id()
+ store i64 %r1, i64* @G_r1
%r2 = tail call spir_func i64 @get_local_linear_id()
+ store i64 %r2, i64* @G_r2
%r3 = tail call spir_func i32 @get_work_dim()
+ store i32 %r3, i32* @G_r3
%r4 = tail call spir_func i32 @get_sub_group_size()
+ store i32 %r4, i32* @G_r4
%r5 = tail call spir_func i32 @get_max_sub_group_size()
+ store i32 %r5, i32* @G_r5
%r6 = tail call spir_func i32 @get_num_sub_groups()
+ store i32 %r6, i32* @G_r6
%r7 = tail call spir_func i32 @get_enqueued_num_sub_groups()
+ store i32 %r7, i32* @G_r7
%r8 = tail call spir_func i32 @get_sub_group_id()
+ store i32 %r8, i32* @G_r8
%r9 = tail call spir_func i32 @get_sub_group_local_id()
+ store i32 %r9, i32* @G_r9
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/debug-info/debug-type-pointer.ll b/llvm/test/CodeGen/SPIRV/debug-info/debug-type-pointer.ll
index 3e0d0cc4cd8e2..d260c9f94d4ad 100644
--- a/llvm/test/CodeGen/SPIRV/debug-info/debug-type-pointer.ll
+++ b/llvm/test/CodeGen/SPIRV/debug-info/debug-type-pointer.ll
@@ -126,6 +126,7 @@ define spir_func i32 @test0() !dbg !17 {
%14 = load ptr addrspace(4), ptr %11, align 4, !dbg !65
store ptr addrspace(4) %14, ptr %12, align 4, !dbg !64
#dbg_declare(ptr %13, !66, !DIExpression(DW_OP_constu, 0, DW_OP_swap, DW_OP_xderef), !70)
+ store [8 x i32] zeroinitializer, ptr %13, align 4
ret i32 0, !dbg !71
}
@@ -169,6 +170,7 @@ define spir_func i32 @test1() !dbg !72 {
%14 = load ptr addrspace(4), ptr %11, align 4, !dbg !97
store ptr addrspace(4) %14, ptr %12, align 4, !dbg !96
#dbg_declare(ptr %13, !98, !DIExpression(DW_OP_constu, 0, DW_OP_swap, DW_OP_xderef), !99)
+ store [8 x i32] zeroinitializer, ptr %13, align 4
ret i32 0, !dbg !100
}
diff --git a/llvm/test/CodeGen/SPIRV/event-zero-const.ll b/llvm/test/CodeGen/SPIRV/event-zero-const.ll
index 523d2ad9825f3..2bf8259e78785 100644
--- a/llvm/test/CodeGen/SPIRV/event-zero-const.ll
+++ b/llvm/test/CodeGen/SPIRV/event-zero-const.ll
@@ -12,11 +12,15 @@
; CHECK: OpINotEqual %[[#]] %[[#]] %[[#LongNull]]
; CHECK: OpGroupAsyncCopy %[[#EventTy]] %[[#]] %[[#]] %[[#]] %[[#]] %[[#]] %[[#EventNull]]
+ at G_r1 = global i1 0
+ at G_e1 = global target("spirv.Event") poison
define weak_odr dso_local spir_kernel void @foo(i64 %_arg_i, ptr addrspace(1) %_arg_ptr, ptr addrspace(3) %_arg_local) {
entry:
%r1 = icmp ne i64 %_arg_i, 0
+ store i1 %r1, ptr @G_r1
%e1 = tail call spir_func target("spirv.Event") @__spirv_GroupAsyncCopy(i32 2, ptr addrspace(3) %_arg_local, ptr addrspace(1) %_arg_ptr, i64 1, i64 1, target("spirv.Event") zeroinitializer)
+ store target("spirv.Event") %e1, ptr @G_e1
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_function_pointers/fun-ptr-addrcast.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_function_pointers/fun-ptr-addrcast.ll
index e5736b88b63a3..a9a0d3358f8cc 100644
--- a/llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_function_pointers/fun-ptr-addrcast.ll
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_INTEL_function_pointers/fun-ptr-addrcast.ll
@@ -11,15 +11,22 @@
@G1 = addrspace(1) constant { [3 x ptr addrspace(4)] } { [3 x ptr addrspace(4)] [ptr addrspace(4) null, ptr addrspace(4) addrspacecast (ptr @foo to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr @bar to ptr addrspace(4))] }
@G2 = addrspace(1) constant { [3 x ptr addrspace(4)] } { [3 x ptr addrspace(4)] [ptr addrspace(4) addrspacecast (ptr null to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr @bar to ptr addrspace(4)), ptr addrspace(4) addrspacecast (ptr @foo to ptr addrspace(4))] }
+ at G_r1_foo = global ptr addrspace(4) null
+ at G_r2_foo = global ptr addrspace(4) null
+ at G_r1_bar = global ptr addrspace(4) null
+
define void @foo(ptr addrspace(4) %p) {
entry:
%r1 = addrspacecast ptr @foo to ptr addrspace(4)
+ store ptr addrspace(4) %r1, ptr @G_r1_foo
%r2 = addrspacecast ptr null to ptr addrspace(4)
+ store ptr addrspace(4) %r2, ptr @G_r2_foo
ret void
}
define void @bar(ptr addrspace(4) %p) {
entry:
%r1 = addrspacecast ptr @bar to ptr addrspace(4)
+ store ptr addrspace(4) %r1, ptr @G_r1_bar
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bfloat16/bfloat16.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bfloat16/bfloat16.ll
index 22668e71fb257..92652f1faefc0 100644
--- a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bfloat16/bfloat16.ll
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_bfloat16/bfloat16.ll
@@ -12,11 +12,16 @@
target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
target triple = "spir64-unknown-unknown"
+ at G1 = global bfloat 0.0
+ at G2 = global <2 x bfloat> zeroinitializer
+
define spir_kernel void @test() {
entry:
%addr1 = alloca bfloat
%addr2 = alloca <2 x bfloat>
%data1 = load bfloat, ptr %addr1
%data2 = load <2 x bfloat>, ptr %addr2
+ store bfloat %data1, ptr @G1
+ store <2 x bfloat> %data2, ptr @G2
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_float_controls2/decoration.ll b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_float_controls2/decoration.ll
index d3fe9e43450cd..81497f26f1aef 100644
--- a/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_float_controls2/decoration.ll
+++ b/llvm/test/CodeGen/SPIRV/extensions/SPV_KHR_float_controls2/decoration.ll
@@ -79,6 +79,54 @@
; CHECK: OpDecorate %[[#maxResV]] FPFastMathMode NotNaN|NotInf|NSZ|AllowRecip|AllowContract|AllowReassoc|AllowTransform
; CHECK: OpDecorate %[[#maxCommonResV]] FPFastMathMode NotNaN|NotInf
+ at G_addRes = global float 0.0
+ at G_subRes = global float 0.0
+ at G_mulRes = global float 0.0
+ at G_divRes = global float 0.0
+ at G_remRes = global float 0.0
+ at G_negRes = global float 0.0
+ at G_oeqRes = global i1 0
+ at G_oneRes = global i1 0
+ at G_oltRes = global i1 0
+ at G_ogtRes = global i1 0
+ at G_oleRes = global i1 0
+ at G_ogeRes = global i1 0
+ at G_ordRes = global i1 0
+ at G_ueqRes = global i1 0
+ at G_uneRes = global i1 0
+ at G_ultRes = global i1 0
+ at G_ugtRes = global i1 0
+ at G_uleRes = global i1 0
+ at G_ugeRes = global i1 0
+ at G_unoRes = global i1 0
+ at G_modRes = global float 0.0
+ at G_maxRes = global float 0.0
+ at G_maxCommonRes = global float 0.0
+
+ at G_addResV = global <2 x float> zeroinitializer
+ at G_subResV = global <2 x float> zeroinitializer
+ at G_mulResV = global <2 x float> zeroinitializer
+ at G_divResV = global <2 x float> zeroinitializer
+ at G_remResV = global <2 x float> zeroinitializer
+ at G_negResV = global <2 x float> zeroinitializer
+ at G_oeqResV = global <2 x i1> zeroinitializer
+ at G_oneResV = global <2 x i1> zeroinitializer
+ at G_oltResV = global <2 x i1> zeroinitializer
+ at G_ogtResV = global <2 x i1> zeroinitializer
+ at G_oleResV = global <2 x i1> zeroinitializer
+ at G_ogeResV = global <2 x i1> zeroinitializer
+ at G_ordResV = global <2 x i1> zeroinitializer
+ at G_ueqResV = global <2 x i1> zeroinitializer
+ at G_uneResV = global <2 x i1> zeroinitializer
+ at G_ultResV = global <2 x i1> zeroinitializer
+ at G_ugtResV = global <2 x i1> zeroinitializer
+ at G_uleResV = global <2 x i1> zeroinitializer
+ at G_ugeResV = global <2 x i1> zeroinitializer
+ at G_unoResV = global <2 x i1> zeroinitializer
+ at G_modResV = global <2 x float> zeroinitializer
+ at G_maxResV = global <2 x float> zeroinitializer
+ at G_maxCommonResV = global <2 x float> zeroinitializer
+
; Function Attrs: convergent mustprogress nofree nounwind willreturn memory(none)
declare spir_func float @_Z4fmodff(float, float)
declare dso_local spir_func noundef nofpclass(nan inf) float @_Z16__spirv_ocl_fmaxff(float noundef nofpclass(nan inf), float noundef nofpclass(nan inf)) local_unnamed_addr #1
@@ -91,55 +139,101 @@ declare dso_local spir_func noundef nofpclass(nan inf) <2 x float> @_Z23__spirv_
define weak_odr dso_local spir_kernel void @foo(float %1, float %2) {
entry:
%addRes = fadd float %1, %2
+ store float %addRes, float* @G_addRes
%subRes = fsub nnan float %1, %2
+ store float %subRes, float* @G_subRes
%mulRes = fmul ninf float %1, %2
+ store float %mulRes, float* @G_mulRes
%divRes = fdiv nsz float %1, %2
+ store float %divRes, float* @G_divRes
%remRes = frem arcp float %1, %2
+ store float %remRes, float* @G_remRes
%negRes = fneg fast float %1
+ store float %negRes, float* @G_negRes
%oeqRes = fcmp nnan ninf oeq float %1, %2
+ store i1 %oeqRes, i1* @G_oeqRes
%oneRes = fcmp one float %1, %2, !spirv.Decorations !3
+ store i1 %oneRes, i1* @G_oneRes
%oltRes = fcmp nnan olt float %1, %2, !spirv.Decorations !3
+ store i1 %oltRes, i1* @G_oltRes
%ogtRes = fcmp ninf ogt float %1, %2, !spirv.Decorations !3
+ store i1 %ogtRes, i1* @G_ogtRes
%oleRes = fcmp nsz ole float %1, %2, !spirv.Decorations !3
+ store i1 %oleRes, i1* @G_oleRes
%ogeRes = fcmp arcp oge float %1, %2, !spirv.Decorations !3
+ store i1 %ogeRes, i1* @G_ogeRes
%ordRes = fcmp fast ord float %1, %2, !spirv.Decorations !3
+ store i1 %ordRes, i1* @G_ordRes
%ueqRes = fcmp nnan ninf ueq float %1, %2, !spirv.Decorations !3
+ store i1 %ueqRes, i1* @G_ueqRes
%uneRes = fcmp une float %1, %2, !spirv.Decorations !3
+ store i1 %uneRes, i1* @G_uneRes
%ultRes = fcmp ult float %1, %2, !spirv.Decorations !3
+ store i1 %ultRes, i1* @G_ultRes
%ugtRes = fcmp ugt float %1, %2, !spirv.Decorations !3
+ store i1 %ugtRes, i1* @G_ugtRes
%uleRes = fcmp ule float %1, %2, !spirv.Decorations !3
+ store i1 %uleRes, i1* @G_uleRes
%ugeRes = fcmp uge float %1, %2, !spirv.Decorations !3
+ store i1 %ugeRes, i1* @G_ugeRes
%unoRes = fcmp uno float %1, %2, !spirv.Decorations !3
+ store i1 %unoRes, i1* @G_unoRes
%modRes = call spir_func float @_Z4fmodff(float %1, float %2)
+ store float %modRes, float* @G_modRes
%maxRes = tail call fast spir_func noundef nofpclass(nan inf) float @_Z16__spirv_ocl_fmaxff(float noundef nofpclass(nan inf) %1, float noundef nofpclass(nan inf) %2)
+ store float %maxRes, float* @G_maxRes
%maxCommonRes = tail call spir_func noundef float @_Z23__spirv_ocl_fmax_commonff(float noundef nofpclass(nan inf) %1, float noundef nofpclass(nan inf) %2)
+ store float %maxCommonRes, float* @G_maxCommonRes
ret void
}
define weak_odr dso_local spir_kernel void @fooV(<2 x float> %v1, <2 x float> %v2) {
%addResV = fadd <2 x float> %v1, %v2
+ store <2 x float> %addResV, <2 x float>* @G_addResV
%subResV = fsub nnan <2 x float> %v1, %v2
+ store <2 x float> %subResV, <2 x float>* @G_subResV
%mulResV = fmul ninf <2 x float> %v1, %v2
+ store <2 x float> %mulResV, <2 x float>* @G_mulResV
%divResV = fdiv nsz <2 x float> %v1, %v2
+ store <2 x float> %divResV, <2 x float>* @G_divResV
%remResV = frem arcp <2 x float> %v1, %v2
+ store <2 x float> %remResV, <2 x float>* @G_remResV
%negResV = fneg fast <2 x float> %v1
+ store <2 x float> %negResV, <2 x float>* @G_negResV
%oeqResV = fcmp nnan ninf oeq <2 x float> %v1, %v2
+ store <2 x i1> %oeqResV, <2 x i1>* @G_oeqResV
%oneResV = fcmp one <2 x float> %v1, %v2, !spirv.Decorations !3
+ store <2 x i1> %oneResV, <2 x i1>* @G_oneResV
%oltResV = fcmp nnan olt <2 x float> %v1, %v2, !spirv.Decorations !3
+ store <2 x i1> %oltResV, <2 x i1>* @G_oltResV
%ogtResV = fcmp ninf ogt <2 x float> %v1, %v2, !spirv.Decorations !3
+ store <2 x i1> %ogtResV, <2 x i1>* @G_ogtResV
%oleResV = fcmp nsz ole <2 x float> %v1, %v2, !spirv.Decorations !3
+ store <2 x i1> %oleResV, <2 x i1>* @G_oleResV
%ogeResV = fcmp arcp oge <2 x float> %v1, %v2, !spirv.Decorations !3
+ store <2 x i1> %ogeResV, <2 x i1>* @G_ogeResV
%ordResV = fcmp fast ord <2 x float> %v1, %v2, !spirv.Decorations !3
+ store <2 x i1> %ordResV, <2 x i1>* @G_ordResV
%ueqResV = fcmp nnan ninf ueq <2 x float> %v1, %v2, !spirv.Decorations !3
+ store <2 x i1> %ueqResV, <2 x i1>* @G_ueqResV
%uneResV = fcmp une <2 x float> %v1, %v2, !spirv.Decorations !3
+ store <2 x i1> %uneResV, <2 x i1>* @G_uneResV
%ultResV = fcmp ult <2 x float> %v1, %v2, !spirv.Decorations !3
+ store <2 x i1> %ultResV, <2 x i1>* @G_ultResV
%ugtResV = fcmp ugt <2 x float> %v1, %v2, !spirv.Decorations !3
+ store <2 x i1> %ugtResV, <2 x i1>* @G_ugtResV
%uleResV = fcmp ule <2 x float> %v1, %v2, !spirv.Decorations !3
+ store <2 x i1> %uleResV, <2 x i1>* @G_uleResV
%ugeResV = fcmp uge <2 x float> %v1, %v2, !spirv.Decorations !3
+ store <2 x i1> %ugeResV, <2 x i1>* @G_ugeResV
%unoResV = fcmp uno <2 x float> %v1, %v2, !spirv.Decorations !3
+ store <2 x i1> %unoResV, <2 x i1>* @G_unoResV
%modResV = call spir_func <2 x float> @_Z4fmodDv2_fDv2_f(<2 x float> %v1, <2 x float> %v2)
+ store <2 x float> %modResV, <2 x float>* @G_modResV
%maxResV = tail call fast spir_func noundef nofpclass(nan inf) <2 x float> @_Z16__spirv_ocl_fmaxDv2_fDv2_f(<2 x float> noundef nofpclass(nan inf) %v1, <2 x float> noundef nofpclass(nan inf) %v2)
+ store <2 x float> %maxResV, <2 x float>* @G_maxResV
%maxCommonResV = tail call spir_func noundef <2 x float> @_Z23__spirv_ocl_fmax_commonDv2_fDv2_f(<2 x float> noundef nofpclass(nan inf) %v1, <2 x float> noundef nofpclass(nan inf) %v2)
+ store <2 x float> %maxCommonResV, <2 x float>* @G_maxCommonResV
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/extensions/enable-all-extensions-but-one.ll b/llvm/test/CodeGen/SPIRV/extensions/enable-all-extensions-but-one.ll
index 4db0ba33d52c9..face4a9f5e615 100644
--- a/llvm/test/CodeGen/SPIRV/extensions/enable-all-extensions-but-one.ll
+++ b/llvm/test/CodeGen/SPIRV/extensions/enable-all-extensions-but-one.ll
@@ -2,10 +2,15 @@
; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv32-unknown-unknown --spirv-ext=KHR %s -o - | FileCheck %s
; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv32-unknown-unknown --spirv-ext=khr %s -o - | FileCheck %s
+ at G = global i32 0
+
define i6 @foo() {
%call = tail call i32 @llvm.bitreverse.i32(i32 42)
+ store i32 %call, ptr @G
ret i6 2
}
; CHECK-NOT: OpExtension "SPV_INTEL_arbitrary_precision_integers"
; CHECK-DAG: OpExtension "SPV_KHR_bit_instructions"
+
+declare i32 @llvm.bitreverse.i32(i32)
diff --git a/llvm/test/CodeGen/SPIRV/freeze.ll b/llvm/test/CodeGen/SPIRV/freeze.ll
index 9077d2ede72a9..e8c9e933904cc 100644
--- a/llvm/test/CodeGen/SPIRV/freeze.ll
+++ b/llvm/test/CodeGen/SPIRV/freeze.ll
@@ -1,15 +1,15 @@
; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
; TODO: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %}
-; CHECK: OpName %[[Arg1:.*]] "arg1"
-; CHECK: OpName %[[Arg2:.*]] "arg2"
-; CHECK: OpName %[[NotAStaticPoison:.*]] "poison1"
-; CHECK: OpName %[[NotAStaticPoison]] "nil0"
-; CHECK: OpName %[[StaticPoisonIntFreeze:.*]] "nil1"
-; CHECK: OpName %[[StaticPoisonFloatFreeze:.*]] "nil2"
-; CHECK: OpName %[[Arg1]] "val1"
-; CHECK: OpName %[[Const100:.*]] "val2"
-; CHECK: OpName %[[Const100]] "val3"
+; CHECK-DAG: OpName %[[Arg1:.*]] "arg1"
+; CHECK-DAG: OpName %[[Arg2:.*]] "arg2"
+; CHECK-DAG: OpName %[[NotAStaticPoison:.*]] "poison1"
+; CHECK-DAG: OpName %[[NotAStaticPoison]] "nil0"
+; CHECK-DAG: OpName %[[StaticPoisonIntFreeze:.*]] "nil1"
+; CHECK-DAG: OpName %[[StaticPoisonFloatFreeze:.*]] "nil2"
+; CHECK-DAG: OpName %[[Arg1]] "val1"
+; CHECK-DAG: OpName %[[Const100:.*]] "val2"
+; CHECK-DAG: OpName %[[Const100]] "val3"
; CHECK: OpDecorate
; CHECK-DAG: %[[FloatTy:.*]] = OpTypeFloat 32
; CHECK-DAG: %[[ShortTy:.*]] = OpTypeInt 16 0
@@ -18,17 +18,37 @@
; CHECK-DAG: %[[Undef32:.*]] = OpUndef %[[IntTy]]
; CHECK-DAG: %[[UndefFloat:.*]] = OpUndef %[[FloatTy]]
; CHECK-DAG: %[[Const100]] = OpConstant %[[IntTy]] 100
-; CHECK: %[[Arg1]] = OpFunctionParameter %[[FloatTy]]
-; CHECK: %[[NotAStaticPoison]] = OpIAdd %[[ShortTy]] %[[Arg2]] %[[Undef16]]
-define spir_func void @foo(float %arg1, i16 %arg2) {
+define spir_func i16 @test_nil0(i16 %arg2) {
entry:
+; CHECK: %[[NotAStaticPoison]] = OpIAdd %[[ShortTy]] %[[Arg2]] %[[Undef16]]
%poison1 = add i16 %arg2, undef
%nil0 = freeze i16 %poison1
+ ret i16 %nil0
+}
+
+define spir_func i32 @test_nil1() {
+entry:
%nil1 = freeze i32 undef
+ ret i32 %nil1
+}
+
+define spir_func float @test_nil2() {
+entry:
%nil2 = freeze float poison
+ ret float %nil2
+}
+
+define spir_func float @freeze_float(float %arg1) {
+entry:
+; CHECK: %[[Arg1]] = OpFunctionParameter %[[FloatTy]]
%val1 = freeze float %arg1
+ ret float %val1
+}
+
+define spir_func i32 @foo() {
+entry:
%val2 = freeze i32 100
%val3 = freeze i32 %val2
- ret void
-}
+ ret i32 %val3
+}
\ No newline at end of file
diff --git a/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/AddUint64.ll b/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/AddUint64.ll
index a97492b8453ea..a15d628cc3614 100644
--- a/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/AddUint64.ll
+++ b/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/AddUint64.ll
@@ -63,7 +63,7 @@ entry:
; CHECK: %[[#a_high:]] = OpVectorShuffle %[[#vec2_int_32]] %[[#a]] %[[#undef_v4i32]] 1 3
; CHECK: %[[#b_low:]] = OpVectorShuffle %[[#vec2_int_32]] %[[#b]] %[[#undef_v4i32]] 0 2
; CHECK: %[[#b_high:]] = OpVectorShuffle %[[#vec2_int_32]] %[[#b]] %[[#undef_v4i32]] 1 3
-; CHECK: %[[#iaddcarry:]] = OpIAddCarry %[[#struct_v2i32_v2i32]] %[[#a_low]] %[[#vec2_int_32]]
+; CHECK: %[[#iaddcarry:]] = OpIAddCarry %[[#struct_v2i32_v2i32]] %[[#a_low]] %[[#b_low]]
; CHECK: %[[#lowsum:]] = OpCompositeExtract %[[#vec2_int_32]] %[[#iaddcarry]] 0
; CHECK: %[[#carry:]] = OpCompositeExtract %[[#vec2_int_32]] %[[#iaddcarry]] 1
; CHECK: %[[#carry_ne0:]] = OpINotEqual %[[#vec2_bool]] %[[#carry]] %[[#const_v2i32_0_0]]
diff --git a/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/abs.ll b/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/abs.ll
index 4a15fa8b14537..75fac211f1108 100644
--- a/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/abs.ll
+++ b/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/abs.ll
@@ -3,24 +3,25 @@
; CHECK: OpExtInstImport "GLSL.std.450"
+ at i = global i32 0, align 4
+ at absi = global i32 0, align 4
+ at f = global float 0.0, align 4
+ at absf = global float 0.0, align 4
+
define void @main() #1 {
entry:
- %i = alloca i32, align 4
- %absi = alloca i32, align 4
- %f = alloca float, align 4
- %absf = alloca float, align 4
- %0 = load i32, ptr %i, align 4
+ %0 = load i32, ptr @i, align 4
; CHECK: %[[#]] = OpExtInst %[[#]] %[[#]] SAbs %[[#]]
%elt.abs = call i32 @llvm.abs.i32(i32 %0, i1 false)
- store i32 %elt.abs, ptr %absi, align 4
- %1 = load float, ptr %f, align 4
+ store i32 %elt.abs, ptr @absi, align 4
+ %1 = load float, ptr @f, align 4
; CHECK: %[[#]] = OpExtInst %[[#]] %[[#]] FAbs %[[#]]
%elt.abs1 = call float @llvm.fabs.f32(float %1)
- store float %elt.abs1, ptr %absf, align 4
+ store float %elt.abs1, ptr @absf, align 4
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/log10.ll b/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/log10.ll
index 7583066c01cf8..dceaa8c209957 100644
--- a/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/log10.ll
+++ b/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/log10.ll
@@ -7,21 +7,23 @@
; CHECK: %[[#v4float:]] = OpTypeVector %[[#float]] 4
; CHECK: %[[#float_0_30103001:]] = OpConstant %[[#float]] 0.30103000998497009
+ at logf = global float 0.0, align 4
+ at logf4 = global <4 x float> zeroinitializer, align 16
+
define void @main(float %f, <4 x float> %f4) {
entry:
; CHECK-DAG: %[[#f:]] = OpFunctionParameter %[[#float]]
; CHECK-DAG: %[[#f4:]] = OpFunctionParameter %[[#v4float]]
- %logf = alloca float, align 4
- %logf4 = alloca <4 x float>, align 16
-
; CHECK: %[[#log2:]] = OpExtInst %[[#float]] %[[#extinst]] Log2 %[[#f]]
; CHECK: %[[#res:]] = OpFMul %[[#float]] %[[#log2]] %[[#float_0_30103001]]
%elt.log10 = call float @llvm.log10.f32(float %f)
+ store float %elt.log10, ptr @logf, align 4
; CHECK: %[[#log2:]] = OpExtInst %[[#v4float]] %[[#extinst]] Log2 %[[#f4]]
; CHECK: %[[#res:]] = OpVectorTimesScalar %[[#v4float]] %[[#log2]] %[[#float_0_30103001]]
%elt.log101 = call <4 x float> @llvm.log10.v4f32(<4 x float> %f4)
+ store <4 x float> %elt.log101, ptr @logf4, align 16
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/instructions/insertvalue-undef-ptr.ll b/llvm/test/CodeGen/SPIRV/instructions/insertvalue-undef-ptr.ll
index b788f34bf7238..02825e3cbb599 100644
--- a/llvm/test/CodeGen/SPIRV/instructions/insertvalue-undef-ptr.ll
+++ b/llvm/test/CodeGen/SPIRV/instructions/insertvalue-undef-ptr.ll
@@ -4,25 +4,40 @@
; CHECK-LABEL: Begin function original_testcase
define fastcc void @original_testcase() {
top:
+ %0 = alloca [1 x ptr], align 4
; CHECK: OpCompositeInsert
- %0 = insertvalue [1 x ptr] zeroinitializer, ptr poison, 0
+ %1 = insertvalue [1 x ptr] zeroinitializer, ptr poison, 0
+ store [1 x ptr] %1, ptr %0
ret void
}
; CHECK-LABEL: Begin function additional_testcases
define fastcc void @additional_testcases() {
top:
+ %0 = alloca [2 x ptr], align 4
+
+
; Test with different pointer types
; CHECK: OpCompositeInsert
%1 = insertvalue [1 x ptr] zeroinitializer, ptr undef, 0
+ ; CHECK: OpStore
+ store [1 x ptr] %1, ptr %0
+
; CHECK-NEXT: OpCompositeInsert
%2 = insertvalue {ptr, i32} zeroinitializer, ptr poison, 0
+ ; CHECK: OpStore
+ store {ptr, i32} %2, ptr %0
+
; CHECK-NEXT: OpCompositeInsert
%3 = insertvalue {ptr, ptr} undef, ptr null, 0
+ ; CHECK: OpStore
+ store {ptr, ptr} %3, ptr %0
; Test with undef aggregate
; CHECK-NEXT: OpCompositeInsert
%4 = insertvalue [1 x ptr] undef, ptr undef, 0
+ ; CHECK: OpStore
+ store [1 x ptr] %4, ptr %0
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/instructions/select-ptr-load.ll b/llvm/test/CodeGen/SPIRV/instructions/select-ptr-load.ll
index 6e6cd2f68a971..510c7954c78f8 100644
--- a/llvm/test/CodeGen/SPIRV/instructions/select-ptr-load.ll
+++ b/llvm/test/CodeGen/SPIRV/instructions/select-ptr-load.ll
@@ -13,13 +13,18 @@
%struct = type { [3 x float] }
+ at G = global float 0.0
+
define spir_kernel void @bar(i1 %sw) {
entry:
%var1 = alloca %struct
+ store %struct zeroinitializer, ptr %var1
%var2 = alloca %struct
+ store %struct zeroinitializer, ptr %var2
%elem1 = getelementptr inbounds [3 x float], ptr %var1, i64 0, i64 0
%elem2 = getelementptr inbounds [3 x float], ptr %var2, i64 0, i64 1
%elem = select i1 %sw, ptr %elem1, ptr %elem2
%res = load float, ptr %elem
+ store float %res, ptr @G
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/keep-tracked-const.ll b/llvm/test/CodeGen/SPIRV/keep-tracked-const.ll
deleted file mode 100644
index efde6a2c082fc..0000000000000
--- a/llvm/test/CodeGen/SPIRV/keep-tracked-const.ll
+++ /dev/null
@@ -1,23 +0,0 @@
-; This test case ensures that cleaning of temporary constants doesn't purge tracked ones.
-
-; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
-; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
-
-; CHECK-SPIRV-DAG: %[[#Int:]] = OpTypeInt 8 0
-; CHECK-SPIRV-DAG: %[[#C0:]] = OpConstantNull %[[#Int]]
-; CHECK-SPIRV-DAG: %[[#C1:]] = OpConstant %[[#Int]] 1{{$}}
-
-define spir_kernel void @foo() {
-entry:
- %addr = alloca i32
- %r1 = call i8 @_Z20__spirv_SpecConstantia(i32 0, i8 1)
- ; The name '%conv17.i' is important for the test case,
- ; because it includes i32 0 when encoded for SPIR-V usage.
- %conv17.i = sext i8 %r1 to i64
- %tobool = trunc i8 %r1 to i1
- %r2 = zext i1 %tobool to i32
- store i32 %r2, ptr %addr
- ret void
-}
-
-declare i8 @_Z20__spirv_SpecConstantia(i32, i8)
diff --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/assume.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/assume.ll
index 3d2080e0050b7..691325251f11d 100644
--- a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/assume.ll
+++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/assume.ll
@@ -8,14 +8,15 @@
%class.anon = type { i8 }
-define spir_func void @_Z3fooi(i32 %x) {
+define spir_func i32 @_Z3fooi(i32 %x) {
entry:
%x.addr = alloca i32, align 4
store i32 %x, i32* %x.addr, align 4
- %0 = load i32, i32* %x.addr, align 4
+ %0 = load i32, ptr %x.addr, align 4
%cmp = icmp ne i32 %0, 0
call void @llvm.assume(i1 %cmp)
- ret void
+ %retval = select i1 %cmp, i32 100, i32 10
+ ret i32 %retval
}
declare void @llvm.assume(i1)
@@ -45,9 +46,9 @@ entry:
call void @llvm.lifetime.start.p0i8(i64 4, i8* %0)
store i32 1, i32* %a, align 4
%1 = load i32, i32* %a, align 4
- call spir_func void @_Z3fooi(i32 %1)
- %2 = bitcast i32* %a to i8*
- call void @llvm.lifetime.end.p0i8(i64 4, i8* %2)
+ %2 = call spir_func i32 @_Z3fooi(i32 %1)
+ %3 = bitcast i32* %a to i8*
+ call void @llvm.lifetime.end.p0i8(i64 4, i8* %3)
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/bitreverse_small_type.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/bitreverse_small_type.ll
index 438fff6e94f89..18856147896bb 100644
--- a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/bitreverse_small_type.ll
+++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/bitreverse_small_type.ll
@@ -7,20 +7,20 @@
; CHECK: OpCapability ArbitraryPrecisionIntegersINTEL
; CHECK: OpExtension "SPV_INTEL_arbitrary_precision_integers"
-; CHECK: %[[#I4:]] = OpTypeInt 4 0
-; CHECK: %[[#I2:]] = OpTypeInt 2 0
-; CHECK: %[[#Z4:]] = OpConstantNull %[[#I4]]
-; CHECK: %[[#Z2:]] = OpConstantNull %[[#I2]]
-; CHECK: %[[#V2I2:]] = OpTypeVector %[[#I2]] 2
-; CHECK: %[[#V2I4:]] = OpTypeVector %[[#I4]] 2
-; CHECK: %[[#V3I2:]] = OpTypeVector %[[#I2]] 3
-; CHECK: %[[#V3I4:]] = OpTypeVector %[[#I4]] 3
-; CHECK: %[[#V4I2:]] = OpTypeVector %[[#I2]] 4
-; CHECK: %[[#V4I4:]] = OpTypeVector %[[#I4]] 4
-; CHECK: %[[#V8I2:]] = OpTypeVector %[[#I2]] 8
-; CHECK: %[[#V8I4:]] = OpTypeVector %[[#I4]] 8
-; CHECK: %[[#V16I2:]] = OpTypeVector %[[#I2]] 16
-; CHECK: %[[#V16I4:]] = OpTypeVector %[[#I4]] 16
+; CHECK-DAG: %[[#I4:]] = OpTypeInt 4 0
+; CHECK-DAG: %[[#I2:]] = OpTypeInt 2 0
+; CHECK-DAG: %[[#Z4:]] = OpConstantNull %[[#I4]]
+; CHECK-DAG: %[[#Z2:]] = OpConstantNull %[[#I2]]
+; CHECK-DAG: %[[#V2I2:]] = OpTypeVector %[[#I2]] 2
+; CHECK-DAG: %[[#V2I4:]] = OpTypeVector %[[#I4]] 2
+; CHECK-DAG: %[[#V3I2:]] = OpTypeVector %[[#I2]] 3
+; CHECK-DAG: %[[#V3I4:]] = OpTypeVector %[[#I4]] 3
+; CHECK-DAG: %[[#V4I2:]] = OpTypeVector %[[#I2]] 4
+; CHECK-DAG: %[[#V4I4:]] = OpTypeVector %[[#I4]] 4
+; CHECK-DAG: %[[#V8I2:]] = OpTypeVector %[[#I2]] 8
+; CHECK-DAG: %[[#V8I4:]] = OpTypeVector %[[#I4]] 8
+; CHECK-DAG: %[[#V16I2:]] = OpTypeVector %[[#I2]] 16
+; CHECK-DAG: %[[#V16I4:]] = OpTypeVector %[[#I4]] 16
; CHECK: %[[#]] = OpBitReverse %[[#I2]] %[[#Z2]]
@@ -36,45 +36,70 @@
; CHECK: %[[#]] = OpBitReverse %[[#V16I2]] %[[#]]
; CHECK: %[[#]] = OpBitReverse %[[#V16I4]] %[[#]]
+ at G_i2_res = global i2 0
+ at G_i4_res = global i4 0
+ at G_v2i2_res = global <2 x i2> zeroinitializer
+ at G_v2i4_res = global <2 x i4> zeroinitializer
+ at G_v3i2_res = global <3 x i2> zeroinitializer
+ at G_v3i4_res = global <3 x i4> zeroinitializer
+ at G_v4i2_res = global <4 x i2> zeroinitializer
+ at G_v4i4_res = global <4 x i4> zeroinitializer
+ at G_v8i2_res = global <8 x i2> zeroinitializer
+ at G_v8i4_res = global <8 x i4> zeroinitializer
+ at G_v16i2_res = global <16 x i2> zeroinitializer
+ at G_v16i4_res = global <16 x i4> zeroinitializer
+
define spir_kernel void @testBitRev() {
entry:
%call2 = call i2 @llvm.bitreverse.i2(i2 0)
+ store i2 %call2, i2* @G_i2_res
%call4 = call i4 @llvm.bitreverse.i4(i4 0)
+ store i4 %call4, i4* @G_i4_res
ret void
}
define spir_kernel void @testBitRevV2(<2 x i2> %a, <2 x i4> %b) {
entry:
%call2 = call <2 x i2> @llvm.bitreverse.v2i2(<2 x i2> %a)
+ store <2 x i2> %call2, <2 x i2>* @G_v2i2_res
%call4 = call <2 x i4> @llvm.bitreverse.v2i4(<2 x i4> %b)
+ store <2 x i4> %call4, <2 x i4>* @G_v2i4_res
ret void
}
define spir_kernel void @testBitRevV3(<3 x i2> %a, <3 x i4> %b) {
entry:
%call2 = call <3 x i2> @llvm.bitreverse.v3i2(<3 x i2> %a)
+ store <3 x i2> %call2, <3 x i2>* @G_v3i2_res
%call4 = call <3 x i4> @llvm.bitreverse.v3i4(<3 x i4> %b)
+ store <3 x i4> %call4, <3 x i4>* @G_v3i4_res
ret void
}
define spir_kernel void @testBitRevV4(<4 x i2> %a, <4 x i4> %b) {
entry:
%call2 = call <4 x i2> @llvm.bitreverse.v4i2(<4 x i2> %a)
+ store <4 x i2> %call2, <4 x i2>* @G_v4i2_res
%call4 = call <4 x i4> @llvm.bitreverse.v4i4(<4 x i4> %b)
+ store <4 x i4> %call4, <4 x i4>* @G_v4i4_res
ret void
}
define spir_kernel void @testBitRevV8(<8 x i2> %a, <8 x i4> %b) {
entry:
%call2 = call <8 x i2> @llvm.bitreverse.v8i2(<8 x i2> %a)
+ store <8 x i2> %call2, <8 x i2>* @G_v8i2_res
%call4 = call <8 x i4> @llvm.bitreverse.v8i4(<8 x i4> %b)
+ store <8 x i4> %call4, <8 x i4>* @G_v8i4_res
ret void
}
define spir_kernel void @testBitRevV16(<16 x i2> %a, <16 x i4> %b) {
entry:
%call2 = call <16 x i2> @llvm.bitreverse.v16i2(<16 x i2> %a)
+ store <16 x i2> %call2, <16 x i2>* @G_v16i2_res
%call4 = call <16 x i4> @llvm.bitreverse.v16i4(<16 x i4> %b)
+ store <16 x i4> %call4, <16 x i4>* @G_v16i4_res
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/constrained-arithmetic.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/constrained-arithmetic.ll
index 11bedfa605f9b..8e8e4df8fabc6 100644
--- a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/constrained-arithmetic.ll
+++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/constrained-arithmetic.ll
@@ -23,15 +23,28 @@
; CHECK: OpExtInst %[[#]] %[[#]] fma %[[#]] %[[#]] %[[#]]
; CHECK: OpFRem
+ at G_r1 = global float 0.0
+ at G_r2 = global float 0.0
+ at G_r3 = global float 0.0
+ at G_r4 = global float 0.0
+ at G_r5 = global float 0.0
+ at G_r6 = global float 0.0
+
; Function Attrs: norecurse nounwind strictfp
define dso_local spir_kernel void @test(float %a, i32 %in, i32 %ui) {
entry:
%r1 = tail call float @llvm.experimental.constrained.fadd.f32(float %a, float %a, metadata !"round.tonearest", metadata !"fpexcept.strict")
+ store float %r1, ptr @G_r1
%r2 = tail call float @llvm.experimental.constrained.fdiv.f32(float %a, float %a, metadata !"round.towardzero", metadata !"fpexcept.strict")
+ store float %r2, ptr @G_r2
%r3 = tail call float @llvm.experimental.constrained.fsub.f32(float %a, float %a, metadata !"round.upward", metadata !"fpexcept.strict")
+ store float %r3, ptr @G_r3
%r4 = tail call float @llvm.experimental.constrained.fmul.f32(float %a, float %a, metadata !"round.downward", metadata !"fpexcept.strict")
+ store float %r4, ptr @G_r4
%r5 = tail call float @llvm.experimental.constrained.fma.f32(float %a, float %a, float %a, metadata !"round.dynamic", metadata !"fpexcept.strict")
+ store float %r5, ptr @G_r5
%r6 = tail call float @llvm.experimental.constrained.frem.f32(float %a, float %a, metadata !"round.dynamic", metadata !"fpexcept.strict")
+ store float %r6, ptr @G_r6
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/lifetime.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/lifetime.ll
index f83cd8ad1969c..375da5b32e232 100644
--- a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/lifetime.ll
+++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/lifetime.ll
@@ -18,19 +18,20 @@
; CL: %[[#FooVar:]] = OpVariable
; CL-NEXT: %[[#Casted1:]] = OpBitcast %[[#PtrChar]] %[[#FooVar]]
; CL-NEXT: OpLifetimeStart %[[#Casted1]] 16
-; CL-NEXT: OpBitcast
-; CL-NEXT: OpInBoundsPtrAccessChain
-; CL-NEXT: %[[#Casted2:]] = OpBitcast %[[#PtrChar]] %[[#FooVar]]
+; CL: OpInBoundsPtrAccessChain
+; CL: %[[#Casted2:]] = OpBitcast %[[#PtrChar]] %[[#FooVar]]
; CL-NEXT: OpLifetimeStop %[[#Casted2]] 16
; VK: OpFunction
; VK: %[[#FooVar:]] = OpVariable
; VK-NEXT: OpInBoundsAccessChain
+; VK-NEXT: OpStore
; VK-NEXT: OpReturn
define spir_func void @foo(ptr noundef byval(%tprange) align 8 %_arg_UserRange) {
%RoundedRangeKernel = alloca %tprange, align 8
call void @llvm.lifetime.start.p0(ptr nonnull %RoundedRangeKernel)
%KernelFunc = getelementptr inbounds i8, ptr %RoundedRangeKernel, i64 8
+ store i64 zeroinitializer, ptr %KernelFunc, align 8
call void @llvm.lifetime.end.p0(ptr nonnull %RoundedRangeKernel)
ret void
}
@@ -39,37 +40,40 @@ define spir_func void @foo(ptr noundef byval(%tprange) align 8 %_arg_UserRange)
; CL: %[[#BarVar:]] = OpVariable
; CL-NEXT: %[[#Casted1:]] = OpBitcast %[[#PtrChar]] %[[#BarVar]]
; CL-NEXT: OpLifetimeStart %[[#Casted1]] 16
-; CL-NEXT: OpBitcast
-; CL-NEXT: OpInBoundsPtrAccessChain
-; CL-NEXT: %[[#Casted2:]] = OpBitcast %[[#PtrChar]] %[[#BarVar]]
+; CL: OpInBoundsPtrAccessChain
+; CL: %[[#Casted2:]] = OpBitcast %[[#PtrChar]] %[[#BarVar]]
; CL-NEXT: OpLifetimeStop %[[#Casted2]] 16
; VK: OpFunction
; VK: %[[#BarVar:]] = OpVariable
; VK-NEXT: OpInBoundsAccessChain
+; VK-NEXT: OpStore
; VK-NEXT: OpReturn
define spir_func void @bar(ptr noundef byval(%tprange) align 8 %_arg_UserRange) {
%RoundedRangeKernel = alloca %tprange, align 8
call void @llvm.lifetime.start.p0(ptr nonnull %RoundedRangeKernel)
%KernelFunc = getelementptr inbounds i8, ptr %RoundedRangeKernel, i64 8
+ store i64 zeroinitializer, ptr %KernelFunc, align 8
call void @llvm.lifetime.end.p0(ptr nonnull %RoundedRangeKernel)
ret void
}
; CL: OpFunction
; CL: %[[#TestVar:]] = OpVariable
-; CL-NEXT: OpLifetimeStart %[[#TestVar]] 1
-; CL-NEXT: OpInBoundsPtrAccessChain
-; CL-NEXT: OpLifetimeStop %[[#TestVar]] 1
+; CL: OpLifetimeStart %[[#TestVar]] 1
+; CL: OpInBoundsPtrAccessChain
+; CL: OpLifetimeStop %[[#TestVar]] 1
; VK: OpFunction
; VK: %[[#Test:]] = OpVariable
; VK-NEXT: OpInBoundsAccessChain
+; VK-NEXT: OpStore
; VK-NEXT: OpReturn
define spir_func void @test(ptr noundef align 8 %_arg) {
%var = alloca i8, align 8
call void @llvm.lifetime.start.p0(ptr nonnull %var)
%KernelFunc = getelementptr inbounds i8, ptr %var, i64 1
+ store i8 0, ptr %KernelFunc, align 8
call void @llvm.lifetime.end.p0(ptr nonnull %var)
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/satur-arith.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/satur-arith.ll
index 08f15c077fed9..db930d1b28ec3 100644
--- a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/satur-arith.ll
+++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/satur-arith.ll
@@ -9,29 +9,55 @@
; CHECK-DAG: OpName %[[#Bar:]] "bar"
; CHECK: %[[#Foo]] = OpFunction
; CHECK: %[[#]] = OpExtInst %[[#]] %[[#]] u_add_sat
-; CHECK-NEXT: %[[#]] = OpExtInst %[[#]] %[[#]] u_sub_sat
-; CHECK-NEXT: %[[#]] = OpExtInst %[[#]] %[[#]] s_add_sat
-; CHECK-NEXT: %[[#]] = OpExtInst %[[#]] %[[#]] s_sub_sat
+; CHECK: %[[#]] = OpExtInst %[[#]] %[[#]] u_sub_sat
+; CHECK: %[[#]] = OpExtInst %[[#]] %[[#]] s_add_sat
+; CHECK: %[[#]] = OpExtInst %[[#]] %[[#]] s_sub_sat
; CHECK: %[[#Bar]] = OpFunction
; CHECK: %[[#]] = OpExtInst %[[#]] %[[#]] u_add_sat
-; CHECK-NEXT: %[[#]] = OpExtInst %[[#]] %[[#]] u_sub_sat
-; CHECK-NEXT: %[[#]] = OpExtInst %[[#]] %[[#]] s_add_sat
-; CHECK-NEXT: %[[#]] = OpExtInst %[[#]] %[[#]] s_sub_sat
+; CHECK: %[[#]] = OpExtInst %[[#]] %[[#]] u_sub_sat
+; CHECK: %[[#]] = OpExtInst %[[#]] %[[#]] s_add_sat
+; CHECK: %[[#]] = OpExtInst %[[#]] %[[#]] s_sub_sat
+
+ at G_r1_foo = global i16 0
+ at G_r2_foo = global i16 0
+ at G_r3_foo = global i16 0
+ at G_r4_foo = global i16 0
+ at G_r1_bar = global <4 x i32> zeroinitializer
+ at G_r2_bar = global <4 x i32> zeroinitializer
+ at G_r3_bar = global <4 x i32> zeroinitializer
+ at G_r4_bar = global <4 x i32> zeroinitializer
define spir_func void @foo(i16 %x, i16 %y) {
entry:
%r1 = tail call i16 @llvm.uadd.sat.i16(i16 %x, i16 %y)
+ store i16 %r1, ptr @G_r1_foo
%r2 = tail call i16 @llvm.usub.sat.i16(i16 %x, i16 %y)
+ store i16 %r2, ptr @G_r2_foo
%r3 = tail call i16 @llvm.sadd.sat.i16(i16 %x, i16 %y)
+ store i16 %r3, ptr @G_r3_foo
%r4 = tail call i16 @llvm.ssub.sat.i16(i16 %x, i16 %y)
+ store i16 %r4, ptr @G_r4_foo
ret void
}
define spir_func void @bar(<4 x i32> %x, <4 x i32> %y) {
entry:
%r1 = tail call <4 x i32> @llvm.uadd.sat.v4i32(<4 x i32> %x, <4 x i32> %y)
+ store <4 x i32> %r1, ptr @G_r1_bar
%r2 = tail call <4 x i32> @llvm.usub.sat.v4i32(<4 x i32> %x, <4 x i32> %y)
+ store <4 x i32> %r2, ptr @G_r2_bar
%r3 = tail call <4 x i32> @llvm.sadd.sat.v4i32(<4 x i32> %x, <4 x i32> %y)
+ store <4 x i32> %r3, ptr @G_r3_bar
%r4 = tail call <4 x i32> @llvm.ssub.sat.v4i32(<4 x i32> %x, <4 x i32> %y)
+ store <4 x i32> %r4, ptr @G_r4_bar
ret void
}
+
+declare i16 @llvm.uadd.sat.i16(i16, i16)
+declare i16 @llvm.usub.sat.i16(i16, i16)
+declare i16 @llvm.sadd.sat.i16(i16, i16)
+declare i16 @llvm.ssub.sat.i16(i16, i16)
+declare <4 x i32> @llvm.uadd.sat.v4i32(<4 x i32>, <4 x i32>)
+declare <4 x i32> @llvm.usub.sat.v4i32(<4 x i32>, <4 x i32>)
+declare <4 x i32> @llvm.sadd.sat.v4i32(<4 x i32>, <4 x i32>)
+declare <4 x i32> @llvm.ssub.sat.v4i32(<4 x i32>, <4 x i32>)
diff --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/uadd.with.overflow.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/uadd.with.overflow.ll
index 08e429f36827c..54cb096da8d89 100644
--- a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/uadd.with.overflow.ll
+++ b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/uadd.with.overflow.ll
@@ -90,12 +90,13 @@ define dso_local spir_func void @umulo_v2i64(<2 x i64> %a, <2 x i64> %b, ptr %p)
; CHECK: OpIAddCarry %[[StructLong]]
; CHECK: OpIAddCarry %[[StructLong]]
; CHECK: OpReturn
-define void @foo(i64 %a, i64 %b) {
+define i64 @foo(i64 %a, i64 %b) {
%r1 = call { i64, i1 } @llvm.uadd.with.overflow.i64(i64 %a, i64 %b)
%r2 = call { i64, i1 } @llvm.uadd.with.overflow.i64(i64 %a, i64 %b)
%d1 = extractvalue { i64, i1 } %r1, 0
%d2 = extractvalue { i64, i1 } %r2, 0
- ret void
+ %sum = add i64 %d1, %d2
+ ret i64 %sum
}
declare {i8, i1} @llvm.uadd.with.overflow.i8(i8, i8)
diff --git a/llvm/test/CodeGen/SPIRV/logical-access-chain.ll b/llvm/test/CodeGen/SPIRV/logical-access-chain.ll
index d56678ecfc2c9..e96ebf777c28f 100644
--- a/llvm/test/CodeGen/SPIRV/logical-access-chain.ll
+++ b/llvm/test/CodeGen/SPIRV/logical-access-chain.ll
@@ -2,6 +2,7 @@
; CHECK-DAG: [[uint:%[0-9]+]] = OpTypeInt 32 0
; CHECK-DAG: [[uint2:%[0-9]+]] = OpTypeVector [[uint]] 2
+; CHECK-DAG: [[uint_0:%[0-9]+]] = OpConstant [[uint]] 0
; CHECK-DAG: [[uint_1:%[0-9]+]] = OpConstant [[uint]] 1
; CHECK-DAG: [[ptr_uint:%[0-9]+]] = OpTypePointer Function [[uint]]
; CHECK-DAG: [[ptr_uint2:%[0-9]+]] = OpTypePointer Function [[uint2]]
@@ -12,7 +13,9 @@ entry:
; CHECK: [[var:%[0-9]+]] = OpVariable [[ptr_uint2]] Function
%1 = getelementptr <2 x i32>, ptr %0, i32 0, i32 1
-; CHECK: {{%[0-9]+}} = OpAccessChain [[ptr_uint]] [[var]] [[uint_1]]
+; CHECK: [[gep:%[0-9]+]] = OpAccessChain [[ptr_uint]] [[var]] [[uint_1]]
+ store i32 0, ptr %1
+; CHECK: OpStore [[gep]] [[uint_0]]
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/logical-struct-access.ll b/llvm/test/CodeGen/SPIRV/logical-struct-access.ll
index 66337b1ba2b37..518e011bf0be2 100644
--- a/llvm/test/CodeGen/SPIRV/logical-struct-access.ll
+++ b/llvm/test/CodeGen/SPIRV/logical-struct-access.ll
@@ -1,5 +1,4 @@
-; RUN: llc -O0 -mtriple=spirv-unknown-vulkan1.3-compute %s -o - -print-after-all | FileCheck %s
-; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv-unknown-vulkan1.3-compute %s -o - -filetype=obj | spirv-val %}
+; RUN: llc -O0 -mtriple=spirv-unknown-vulkan1.3-compute %s -o - | FileCheck %s
; CHECK-DAG: [[uint:%[0-9]+]] = OpTypeInt 32 0
@@ -24,35 +23,85 @@
; CHECK-DAG: [[ptr_A:%[0-9]+]] = OpTypePointer Function [[A]]
; CHECK-DAG: [[ptr_B:%[0-9]+]] = OpTypePointer Function [[B]]
-define void @main() #1 {
-entry:
- %0 = alloca %B, align 4
-; CHECK: [[tmp:%[0-9]+]] = OpVariable [[ptr_B]] Function
-
- %1 = getelementptr %B, ptr %0, i32 0, i32 0
+define internal ptr @gep_B_0(ptr %base) {
+; CHECK: [[tmp:%[0-9]+]] = OpFunctionParameter [[ptr_B]]
; CHECK: {{%[0-9]+}} = OpAccessChain [[ptr_A]] [[tmp]] [[uint_0]]
- %2 = getelementptr inbounds %B, ptr %0, i32 0, i32 0
+ %res = getelementptr %B, ptr %base, i32 0, i32 0
+ ret ptr %res
+}
+
+define internal ptr @gep_inbounds_B_0(ptr %base) {
+; CHECK: [[tmp:%[0-9]+]] = OpFunctionParameter [[ptr_B]]
; CHECK: {{%[0-9]+}} = OpInBoundsAccessChain [[ptr_A]] [[tmp]] [[uint_0]]
+ %res = getelementptr inbounds %B, ptr %base, i32 0, i32 0
+ ret ptr %res
+}
- %3 = getelementptr %B, ptr %0, i32 0, i32 1
+define internal ptr @gep_B_1(ptr %base) {
+; CHECK: [[tmp:%[0-9]+]] = OpFunctionParameter [[ptr_B]]
; CHECK: {{%[0-9]+}} = OpAccessChain [[ptr_uint]] [[tmp]] [[uint_1]]
- %4 = getelementptr inbounds %B, ptr %0, i32 0, i32 1
+ %res = getelementptr %B, ptr %base, i32 0, i32 1
+ ret ptr %res
+}
+
+define internal ptr @gep_inbounds_B_1(ptr %base) {
+; CHECK: [[tmp:%[0-9]+]] = OpFunctionParameter [[ptr_B]]
; CHECK: {{%[0-9]+}} = OpInBoundsAccessChain [[ptr_uint]] [[tmp]] [[uint_1]]
+ %res = getelementptr inbounds %B, ptr %base, i32 0, i32 1
+ ret ptr %res
+}
- %5 = getelementptr %B, ptr %0, i32 0, i32 2
+define internal ptr @gep_B_2(ptr %base) {
+; CHECK: [[tmp:%[0-9]+]] = OpFunctionParameter [[ptr_B]]
; CHECK: {{%[0-9]+}} = OpAccessChain [[ptr_A]] [[tmp]] [[uint_2]]
- %6 = getelementptr inbounds %B, ptr %0, i32 0, i32 2
+ %res = getelementptr %B, ptr %base, i32 0, i32 2
+ ret ptr %res
+}
+
+define internal ptr @gep_inbounds_B_2(ptr %base) {
+; CHECK: [[tmp:%[0-9]+]] = OpFunctionParameter [[ptr_B]]
; CHECK: {{%[0-9]+}} = OpInBoundsAccessChain [[ptr_A]] [[tmp]] [[uint_2]]
+ %res = getelementptr inbounds %B, ptr %base, i32 0, i32 2
+ ret ptr %res
+}
- %7 = getelementptr %B, ptr %0, i32 0, i32 2, i32 1
+define internal ptr @gep_B_2_1(ptr %base) {
+; CHECK: [[tmp:%[0-9]+]] = OpFunctionParameter [[ptr_B]]
; CHECK: {{%[0-9]+}} = OpAccessChain [[ptr_uint]] [[tmp]] [[uint_2]] [[uint_1]]
- %8 = getelementptr inbounds %B, ptr %0, i32 0, i32 2, i32 1
+ %res = getelementptr %B, ptr %base, i32 0, i32 2, i32 1
+ ret ptr %res
+}
+
+define internal ptr @gep_inbounds_B_2_1(ptr %base) {
+; CHECK: [[tmp:%[0-9]+]] = OpFunctionParameter [[ptr_B]]
; CHECK: {{%[0-9]+}} = OpInBoundsAccessChain [[ptr_uint]] [[tmp]] [[uint_2]] [[uint_1]]
+ %res = getelementptr inbounds %B, ptr %base, i32 0, i32 2, i32 1
+ ret ptr %res
+}
- %9 = getelementptr %B, ptr %0, i32 0, i32 2
- %10 = getelementptr %A, ptr %9, i32 0, i32 1
+define internal ptr @gep_B_2_A_1(ptr %base) {
+; CHECK: [[tmp:%[0-9]+]] = OpFunctionParameter [[ptr_B]]
; CHECK: [[x:%[0-9]+]] = OpAccessChain [[ptr_A]] [[tmp]] [[uint_2]]
; CHECK: {{%[0-9]+}} = OpAccessChain [[ptr_uint]] [[x]] [[uint_1]]
+ %x = getelementptr %B, ptr %base, i32 0, i32 2
+ %res = getelementptr %A, ptr %x, i32 0, i32 1
+ ret ptr %res
+}
+
+define void @main() #1 {
+entry:
+ %0 = alloca %B, align 4
+; CHECK: [[tmp:%[0-9]+]] = OpVariable [[ptr_B]] Function
+
+ %1 = call ptr @gep_B_0(ptr %0)
+ %2 = call ptr @gep_inbounds_B_0(ptr %0)
+ %3 = call ptr @gep_B_1(ptr %0)
+ %4 = call ptr @gep_inbounds_B_1(ptr %0)
+ %5 = call ptr @gep_B_2(ptr %0)
+ %6 = call ptr @gep_inbounds_B_2(ptr %0)
+ %7 = call ptr @gep_B_2_1(ptr %0)
+ %8 = call ptr @gep_inbounds_B_2_1(ptr %0)
+ %10 = call ptr @gep_B_2_A_1(ptr %0)
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/phi-insert-point.ll b/llvm/test/CodeGen/SPIRV/phi-insert-point.ll
index 70d121cdf4b3a..a34186d491257 100644
--- a/llvm/test/CodeGen/SPIRV/phi-insert-point.ll
+++ b/llvm/test/CodeGen/SPIRV/phi-insert-point.ll
@@ -36,9 +36,18 @@ ok:
br label %exit
exit:
+ store i64 %r1, ptr @g1
+ store i64 %r2, ptr @g2
+ store ptr addrspace(4) %r3, ptr @g3
+ store ptr addrspace(4) %r4, ptr @g4
ret void
}
+ at g1 = internal global i64 0
+ at g2 = internal global i64 0
+ at g3 = internal global ptr addrspace(4) null
+ at g4 = internal global ptr addrspace(4) null
+
define spir_kernel void @bar(i64 %arg_val, i64 %arg_val_def, ptr addrspace(4) byval(%struct) %arg_ptr, ptr addrspace(4) %arg_ptr_def) {
entry:
%fl = icmp eq i64 %arg_val, 0
@@ -55,5 +64,9 @@ ok:
br label %exit
exit:
+ store i64 %r1, ptr @g1
+ store i64 %r2, ptr @g2
+ store ptr addrspace(4) %r3, ptr @g3
+ store ptr addrspace(4) %r4, ptr @g4
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/phi-ptrcast-dominate.ll b/llvm/test/CodeGen/SPIRV/phi-ptrcast-dominate.ll
index bc090ce55fbec..c250ebae12746 100644
--- a/llvm/test/CodeGen/SPIRV/phi-ptrcast-dominate.ll
+++ b/llvm/test/CodeGen/SPIRV/phi-ptrcast-dominate.ll
@@ -20,11 +20,14 @@
; CHECK: %[[#Case1]] = OpFunction
define spir_func void @case1(i1 %b1, i1 %b2, i1 %b3) {
entry:
+ %tmp.1 = alloca i8, align 1
; CHECK: OpBranchConditional %[[#]] %[[#l1:]] %[[#l2:]]
br i1 %b1, label %l1, label %l2
l1:
%str = phi ptr addrspace(1) [ @.str.1, %entry ], [ @.str.2, %l2 ], [ @.str.2, %l3 ]
+ %v1 = load i8, ptr addrspace(1) %str, align 1
+ store i8 %v1, ptr %tmp.1, align 1
br label %exit
; CHECK: %[[#l2]] = OpLabel
@@ -51,11 +54,14 @@ exit:
; CHECK: %[[#Case2]] = OpFunction
define spir_func void @case2(i1 %b1, i1 %b2, i1 %b3, ptr addrspace(1) byval(%struct1) %str1, ptr addrspace(1) byval(%struct2) %str2) {
entry:
+ %tmp.2 = alloca i8, align 1
; CHECK: OpBranchConditional %[[#]] %[[#l1:]] %[[#l2:]]
br i1 %b1, label %l1, label %l2
l1:
%str = phi ptr addrspace(1) [ %str1, %entry ], [ %str2, %l2 ], [ %str2, %l3 ]
+ %v2 = load i8, ptr addrspace(1) %str, align 1
+ store i8 %v2, ptr %tmp.2, align 1
br label %exit
; CHECK: %[[#l2]] = OpLabel
@@ -83,10 +89,13 @@ define spir_func void @case3(i1 %b1, i1 %b2, i1 %b3, ptr addrspace(1) byval(%str
; CHECK: OpBranchConditional %[[#]] %[[#l1:]] %[[#l2:]]
entry:
+ %tmp.3 = alloca i8, align 1
br i1 %b1, label %l1, label %l2
l1:
%str = phi ptr addrspace(1) [ %_arg_str1, %entry ], [ %str2, %l2 ], [ %str3, %l3 ]
+ %v3 = load i8, ptr addrspace(1) %str, align 1
+ store i8 %v3, ptr %tmp.3, align 1
br label %exit
; CHECK: %[[#l2]] = OpLabel
diff --git a/llvm/test/CodeGen/SPIRV/pointers/bitcast-fix-accesschain.ll b/llvm/test/CodeGen/SPIRV/pointers/bitcast-fix-accesschain.ll
index 7db1eed84bf7d..3382987bbd581 100644
--- a/llvm/test/CodeGen/SPIRV/pointers/bitcast-fix-accesschain.ll
+++ b/llvm/test/CodeGen/SPIRV/pointers/bitcast-fix-accesschain.ll
@@ -26,9 +26,13 @@
%struct.S = type { i32 }
%struct.__wrapper_class = type { [7 x %struct.S] }
+ at G_elem = global ptr null
+ at G_data = global i64 0
+
define spir_kernel void @foo1(ptr noundef byval(%struct.__wrapper_class) align 4 %_arg_Arr) {
entry:
%elem = getelementptr inbounds i8, ptr %_arg_Arr, i64 0
+ store ptr %elem, ptr @G_elem
ret void
}
@@ -36,5 +40,6 @@ define spir_kernel void @foo2(ptr noundef byval(%struct.__wrapper_class) align 4
entry:
%elem = getelementptr inbounds %struct.__wrapper_class, ptr %_arg_Arr, i64 0
%data = load i64, ptr %elem
+ store i64 %data, ptr @G_data
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/pointers/bitcast-fix-load.ll b/llvm/test/CodeGen/SPIRV/pointers/bitcast-fix-load.ll
index d6a0071167cef..ed5652a750582 100644
--- a/llvm/test/CodeGen/SPIRV/pointers/bitcast-fix-load.ll
+++ b/llvm/test/CodeGen/SPIRV/pointers/bitcast-fix-load.ll
@@ -14,8 +14,11 @@
%struct.S = type { i32 }
%struct.__wrapper_class = type { [7 x %struct.S] }
+ at G = global i32 0
+
define spir_kernel void @foo(ptr noundef byval(%struct.__wrapper_class) align 4 %_arg_Arr) {
entry:
%val = load i32, ptr %_arg_Arr
+ store i32 %val, ptr @G
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/pointers/gep-types-1.ll b/llvm/test/CodeGen/SPIRV/pointers/gep-types-1.ll
index 0e2730e18bf38..e47aa61a8acd7 100644
--- a/llvm/test/CodeGen/SPIRV/pointers/gep-types-1.ll
+++ b/llvm/test/CodeGen/SPIRV/pointers/gep-types-1.ll
@@ -30,6 +30,8 @@
%"class.std::complex" = type { { double, double } }
%class.anon = type { i32, ptr addrspace(4), [2 x [2 x %"class.std::complex"]] }
+ at G = global ptr addrspace(4) null
+
define weak_odr dso_local spir_kernel void @foo(i32 noundef %_arg_N, ptr addrspace(1) noundef align 8 %_arg_p) {
entry:
%Kernel = alloca %class.anon, align 8
@@ -38,5 +40,6 @@ entry:
%r0 = addrspacecast ptr addrspace(1) %_arg_p to ptr addrspace(4)
store ptr addrspace(4) %r0, ptr %p, align 8
%r3 = load ptr addrspace(4), ptr %p, align 8
+ store ptr addrspace(4) %r3, ptr @G
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/pointers/getelementptr-addressspace.ll b/llvm/test/CodeGen/SPIRV/pointers/getelementptr-addressspace.ll
index 7a09ac973b590..0e397ec51caaa 100644
--- a/llvm/test/CodeGen/SPIRV/pointers/getelementptr-addressspace.ll
+++ b/llvm/test/CodeGen/SPIRV/pointers/getelementptr-addressspace.ll
@@ -7,9 +7,14 @@
; CHECK: %[[#]] = OpInBoundsPtrAccessChain %[[#PTR1]] %[[#]] %[[#]]
; CHECK: %[[#]] = OpInBoundsPtrAccessChain %[[#PTR2]] %[[#]] %[[#]]
+ at G_c = global ptr addrspace(1) null
+ at G_d = global ptr addrspace(2) null
+
define spir_kernel void @foo(ptr addrspace(1) %a, ptr addrspace(2) %b) {
entry:
%c = getelementptr inbounds i8, ptr addrspace(1) %a, i32 1
+ store ptr addrspace(1) %c, ptr @G_c
%d = getelementptr inbounds i8, ptr addrspace(2) %b, i32 2
+ store ptr addrspace(2) %d, ptr @G_d
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/pointers/getelementptr-base-type.ll b/llvm/test/CodeGen/SPIRV/pointers/getelementptr-base-type.ll
index c822dbc5d6c0e..e12a809125248 100644
--- a/llvm/test/CodeGen/SPIRV/pointers/getelementptr-base-type.ll
+++ b/llvm/test/CodeGen/SPIRV/pointers/getelementptr-base-type.ll
@@ -7,9 +7,12 @@
; CHECK: %[[#GEP:]] = OpInBoundsPtrAccessChain %[[#PTR]] %[[#ARG]] %[[#]]
; CHECK: %[[#]] = OpLoad %[[#FLOAT32]] %[[#GEP]] Aligned 4
+ at G = global float 0.0
+
define spir_kernel void @test1(ptr addrspace(1) %arg1) !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_type_qual !4 {
%a = getelementptr inbounds float, ptr addrspace(1) %arg1, i64 1
%b = load float, ptr addrspace(1) %a, align 4
+ store float %b, ptr @G
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/pointers/getelementptr-bitcast-load.ll b/llvm/test/CodeGen/SPIRV/pointers/getelementptr-bitcast-load.ll
index 1d846a35a65aa..859253e5b18d9 100644
--- a/llvm/test/CodeGen/SPIRV/pointers/getelementptr-bitcast-load.ll
+++ b/llvm/test/CodeGen/SPIRV/pointers/getelementptr-bitcast-load.ll
@@ -7,6 +7,9 @@
; CHECK-DAG: %[[#PTR_VEC3:]] = OpTypePointer CrossWorkgroup %[[#VEC3]]
; CHECK-DAG: %[[#PTR_VEC4:]] = OpTypePointer CrossWorkgroup %[[#VEC4]]
+ at G_loadv1 = global <4 x i8> zeroinitializer
+ at G_loadv2 = global <4 x i8> zeroinitializer
+
; CHECK: %[[#AC1:]] = OpInBoundsPtrAccessChain %[[#PTR_VEC3]] %[[#]] %[[#]]
; CHECK: %[[#BC1:]] = OpBitcast %[[#PTR_VEC4]] %[[#AC1]]
; CHECK: %[[#LD1:]] = OpLoad %[[#VEC4]] %[[#BC1]] Aligned 4
@@ -15,6 +18,7 @@
define spir_kernel void @foo(ptr addrspace(1) %a, i64 %b) {
%index = getelementptr inbounds <3 x i8>, ptr addrspace(1) %a, i64 %b
%loadv = load <4 x i8>, ptr addrspace(1) %index, align 4
+ store <4 x i8> %loadv, ptr @G_loadv1
ret void
}
@@ -29,5 +33,6 @@ define spir_kernel void @bar(ptr addrspace(1) %a, i64 %b) {
; from older LLVM IR with typed pointers.
%cast = bitcast ptr addrspace(1) %index to ptr addrspace(1)
%loadv = load <4 x i8>, ptr addrspace(1) %cast, align 4
+ store <4 x i8> %loadv, ptr @G_loadv2
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/pointers/getelementptr-kernel-arg-char.ll b/llvm/test/CodeGen/SPIRV/pointers/getelementptr-kernel-arg-char.ll
index a5e891dae6f11..3ae03edf5200f 100644
--- a/llvm/test/CodeGen/SPIRV/pointers/getelementptr-kernel-arg-char.ll
+++ b/llvm/test/CodeGen/SPIRV/pointers/getelementptr-kernel-arg-char.ll
@@ -7,11 +7,15 @@
; CHECK-DAG: %[[#PTRINT8:]] = OpTypePointer Workgroup %[[#INT8]]
; CHECK-DAG: %[[#CONST:]] = OpConstant %[[#INT64]] 1
+ at G_gep1 = global ptr addrspace(3) null
+ at G_gep2 = global ptr addrspace(3) null
+
; CHECK: %[[#PARAM1:]] = OpFunctionParameter %[[#PTRINT8]]
define spir_kernel void @test1(ptr addrspace(3) %address) {
; CHECK: %[[#]] = OpInBoundsPtrAccessChain %[[#PTRINT8]] %[[#PARAM1]] %[[#CONST]]
%cast = bitcast ptr addrspace(3) %address to ptr addrspace(3)
%gep = getelementptr inbounds i8, ptr addrspace(3) %cast, i64 1
+ store ptr addrspace(3) %gep, ptr @G_gep1
ret void
}
@@ -19,5 +23,6 @@ define spir_kernel void @test1(ptr addrspace(3) %address) {
define spir_kernel void @test2(ptr addrspace(3) %address) {
; CHECK: %[[#]] = OpInBoundsPtrAccessChain %[[#PTRINT8]] %[[#PARAM2]] %[[#CONST]]
%gep = getelementptr inbounds i8, ptr addrspace(3) %address, i64 1
+ store ptr addrspace(3) %gep, ptr @G_gep2
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/pointers/global-addrspacecast.ll b/llvm/test/CodeGen/SPIRV/pointers/global-addrspacecast.ll
index 19451d23c6830..39563aecafec4 100644
--- a/llvm/test/CodeGen/SPIRV/pointers/global-addrspacecast.ll
+++ b/llvm/test/CodeGen/SPIRV/pointers/global-addrspacecast.ll
@@ -7,13 +7,16 @@
; CHECK-DAG: %[[#value:]] = OpConstant %[[#type]] 456
; CHECK-DAG: %[[#var:]] = OpVariable %[[#ptrty]] Private %[[#value]]
+ at G = internal global i32 0
+
define hidden spir_func void @Foo() {
%p = addrspacecast ptr addrspace(10) @PrivInternal to ptr
%v = load i32, ptr %p, align 4
+ store i32 %v, ptr @G
ret void
; CHECK: OpLabel
-; CHECK-NEXT: OpLoad %[[#type]] %[[#var]] Aligned 4
-; CHECK-Next: OpReturn
+; CHECK: OpLoad %[[#type]] %[[#var]] Aligned 4
+; CHECK: OpReturn
}
define void @main() #1 {
diff --git a/llvm/test/CodeGen/SPIRV/pointers/load-addressspace.ll b/llvm/test/CodeGen/SPIRV/pointers/load-addressspace.ll
index b3c68d22f9bdd..681fb70ad706d 100644
--- a/llvm/test/CodeGen/SPIRV/pointers/load-addressspace.ll
+++ b/llvm/test/CodeGen/SPIRV/pointers/load-addressspace.ll
@@ -9,9 +9,14 @@
; CHECK: %[[#]] = OpLoad %[[#INT8]] %[[#FNP1]] Aligned 1
; CHECK: %[[#]] = OpLoad %[[#INT8]] %[[#FNP2]] Aligned 1
+ at G_c = global i8 0
+ at G_d = global i8 0
+
define spir_kernel void @foo(ptr addrspace(1) %a, ptr addrspace(2) %b) {
entry:
%c = load i8, ptr addrspace(1) %a
+ store i8 %c, ptr @G_c
%d = load i8, ptr addrspace(2) %b
+ store i8 %d, ptr @G_d
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/pointers/phi-chain-types.ll b/llvm/test/CodeGen/SPIRV/pointers/phi-chain-types.ll
index a9e79df259c4f..44134f83cfec3 100644
--- a/llvm/test/CodeGen/SPIRV/pointers/phi-chain-types.ll
+++ b/llvm/test/CodeGen/SPIRV/pointers/phi-chain-types.ll
@@ -51,6 +51,7 @@ l1:
l2:
%val2 = phi ptr addrspace(4) [ %p, %l1 ], [ %val3, %l3 ]
%val1 = phi ptr addrspace(4) [ addrspacecast (ptr addrspace(3) @G1 to ptr addrspace(4)), %l1 ], [ %val2, %l3 ]
+ store i16 0, ptr addrspace(4) %val1, align 2
br i1 %f2, label %l3, label %exit
l3:
@@ -75,6 +76,7 @@ l1:
l2:
%val1 = phi ptr addrspace(4) [ addrspacecast (ptr addrspace(3) @G1 to ptr addrspace(4)), %l1 ], [ %val2, %l3 ]
%val2 = phi ptr addrspace(4) [ %p, %l1 ], [ %val3, %l3 ]
+ store i16 0, ptr addrspace(4) %val1, align 2
br i1 %f2, label %l3, label %exit
exit:
diff --git a/llvm/test/CodeGen/SPIRV/pointers/pointer-addrspacecast.ll b/llvm/test/CodeGen/SPIRV/pointers/pointer-addrspacecast.ll
index 4d5549dfab8d9..123daa411810b 100644
--- a/llvm/test/CodeGen/SPIRV/pointers/pointer-addrspacecast.ll
+++ b/llvm/test/CodeGen/SPIRV/pointers/pointer-addrspacecast.ll
@@ -10,6 +10,7 @@
; CHECK-DAG: OpName %[[#func_chain:]] "chain"
@global = internal addrspace(10) global i32 zeroinitializer
+ at G = global i32 0
define void @simple() {
; CHECK: %[[#func_simple]] = OpFunction
@@ -17,6 +18,7 @@ entry:
%ptr = getelementptr i32, ptr addrspace(10) @global, i32 0
%casted = addrspacecast ptr addrspace(10) %ptr to ptr
%val = load i32, ptr %casted
+ store i32 %val, ptr @G
; CHECK: %{{.*}} = OpLoad %[[#uint]] %[[#var]] Aligned 4
ret void
}
@@ -31,6 +33,7 @@ entry:
%e = addrspacecast ptr addrspace(10) %d to ptr
%val = load i32, ptr %e
+ store i32 %val, ptr @G
; CHECK: %{{.*}} = OpLoad %[[#uint]] %[[#var]] Aligned 4
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/pointers/ptr-eq-types.ll b/llvm/test/CodeGen/SPIRV/pointers/ptr-eq-types.ll
index 876cd3c20cf35..80ee36cfe15d2 100644
--- a/llvm/test/CodeGen/SPIRV/pointers/ptr-eq-types.ll
+++ b/llvm/test/CodeGen/SPIRV/pointers/ptr-eq-types.ll
@@ -15,6 +15,9 @@
; CHECK: OpGenericCastToPtr
; CHECK: OpPtrEqual
+ at G_b1 = global i1 0
+ at G_b2 = global i1 0
+
define spir_kernel void @foo(ptr addrspace(3) align 4 %_arg_local, ptr addrspace(1) align 4 %_arg_global) {
entry:
%p1 = getelementptr inbounds i32, ptr addrspace(1) %_arg_global, i64 0
@@ -24,9 +27,12 @@ entry:
%p4 = addrspacecast ptr addrspace(1) %p3 to ptr addrspace(4)
%p5 = tail call spir_func ptr addrspace(3) @_Z40__spirv_GenericCastToPtrExplicit_ToLocalPvi(ptr addrspace(4) %p4, i32 4)
%b1 = icmp eq ptr addrspace(3) %p5, null
+ store i1 %b1, ptr @G_b1
%p6 = getelementptr inbounds i32, ptr addrspace(3) %p5, i64 0
%p7 = tail call spir_func ptr addrspace(3) @_Z40__spirv_GenericCastToPtrExplicit_ToLocalPvi(ptr addrspace(4) %p4, i32 4)
%b2 = icmp eq ptr addrspace(3) %p7, null
+ store i1 %b2, ptr @G_b2
+ store ptr addrspace(3) %p6, ptr addrspace(3) %p2
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/pointers/resource-vector-load-store.ll b/llvm/test/CodeGen/SPIRV/pointers/resource-vector-load-store.ll
index 7548f4757dbe6..6fc03a386d14d 100644
--- a/llvm/test/CodeGen/SPIRV/pointers/resource-vector-load-store.ll
+++ b/llvm/test/CodeGen/SPIRV/pointers/resource-vector-load-store.ll
@@ -4,18 +4,23 @@
@.str = private unnamed_addr constant [7 x i8] c"buffer\00", align 1
+; The i64 values in the extracts will be turned
+; into immidiate values. There should be no 64-bit
+; integers in the module.
+; CHECK-NOT: OpTypeInt 64 0
+
define void @main() "hlsl.shader"="pixel" {
-; CHECK: %24 = OpFunction %2 None %3 ; -- Begin function main
-; CHECK-NEXT: %1 = OpLabel
-; CHECK-NEXT: %25 = OpVariable %13 Function %22
-; CHECK-NEXT: %26 = OpLoad %7 %23
-; CHECK-NEXT: %27 = OpImageRead %5 %26 %15
-; CHECK-NEXT: %28 = OpCompositeExtract %4 %27 0
-; CHECK-NEXT: %29 = OpCompositeExtract %4 %27 1
-; CHECK-NEXT: %30 = OpFAdd %4 %29 %28
-; CHECK-NEXT: %31 = OpCompositeInsert %5 %30 %27 0
-; CHECK-NEXT: %32 = OpLoad %7 %23
-; CHECK-NEXT: OpImageWrite %32 %15 %31
+; CHECK: %[[FUNC:[0-9]+]] = OpFunction %[[VOID:[0-9]+]] None %[[FNTYPE:[0-9]+]] ; -- Begin function main
+; CHECK-NEXT: %[[LABEL:[0-9]+]] = OpLabel
+; CHECK-NEXT: %[[VAR:[0-9]+]] = OpVariable %[[PTR_FN:[a-zA-Z0-9_]+]] Function %[[INIT:[a-zA-Z0-9_]+]]
+; CHECK-NEXT: %[[LOAD1:[0-9]+]] = OpLoad %[[IMG_TYPE:[a-zA-Z0-9_]+]] %[[IMG_VAR:[a-zA-Z0-9_]+]]
+; CHECK-NEXT: %[[READ:[0-9]+]] = OpImageRead %[[VEC4:[a-zA-Z0-9_]+]] %[[LOAD1]] %[[COORD:[a-zA-Z0-9_]+]]
+; CHECK-NEXT: %[[EXTRACT1:[0-9]+]] = OpCompositeExtract %[[FLOAT:[a-zA-Z0-9_]+]] %[[READ]] 0
+; CHECK-NEXT: %[[EXTRACT2:[0-9]+]] = OpCompositeExtract %[[FLOAT]] %[[READ]] 1
+; CHECK-NEXT: %[[ADD:[0-9]+]] = OpFAdd %[[FLOAT]] %[[EXTRACT2]] %[[EXTRACT1]]
+; CHECK-NEXT: %[[INSERT:[0-9]+]] = OpCompositeInsert %[[VEC4]] %[[ADD]] %[[READ]] 0
+; CHECK-NEXT: %[[LOAD2:[0-9]+]] = OpLoad %[[IMG_TYPE]] %[[IMG_VAR]]
+; CHECK-NEXT: OpImageWrite %[[LOAD2]] %[[COORD]] %[[INSERT]]
; CHECK-NEXT: OpReturn
; CHECK-NEXT: OpFunctionEnd
entry:
diff --git a/llvm/test/CodeGen/SPIRV/pointers/type-deduce-call-no-bitcast.ll b/llvm/test/CodeGen/SPIRV/pointers/type-deduce-call-no-bitcast.ll
index 101116f437811..7409b3db51948 100644
--- a/llvm/test/CodeGen/SPIRV/pointers/type-deduce-call-no-bitcast.ll
+++ b/llvm/test/CodeGen/SPIRV/pointers/type-deduce-call-no-bitcast.ll
@@ -34,6 +34,8 @@
%class.CustomType = type { i64 }
+ at G = global ptr addrspace(4) null
+
define linkonce_odr dso_local spir_func void @bar(ptr addrspace(4) noundef %first) {
entry:
%first.addr = alloca ptr addrspace(4)
@@ -44,6 +46,7 @@ entry:
call spir_func void @foo(i64 noundef 100, ptr addrspace(4) noundef dereferenceable(8) %first.addr.ascast, ptr addrspace(4) noundef dereferenceable(8) %temp.ascast)
call spir_func void @foo(i64 noundef 100, ptr addrspace(4) noundef dereferenceable(8) %temp.ascast, ptr addrspace(4) noundef dereferenceable(8) %first.addr.ascast)
%var = alloca ptr addrspace(4), align 8
+ store ptr addrspace(4) null, ptr %var
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/remove-dead-type-intrinsics.ll b/llvm/test/CodeGen/SPIRV/remove-dead-type-intrinsics.ll
new file mode 100644
index 0000000000000..6bd640f813142
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/remove-dead-type-intrinsics.ll
@@ -0,0 +1,31 @@
+; RUN: llc -O0 -mtriple=spirv-unknown-vulkan1.3-compute %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv-unknown-vulkan1.3-compute %s -o - -filetype=obj | spirv-val %}
+
+%A = type {
+ i32,
+ i32
+}
+
+%B = type {
+ %A,
+ i32,
+ %A
+}
+
+; Make sure all struct types are removed.
+; CHECK-NOT: OpTypeStruct
+
+; Make sure the GEPs and the function scope variable are removed.
+; CHECK: OpFunction
+; CHECK-NEXT: OpLabel
+; CHECK-NEXT: OpReturn
+; CHECK-NEXT: OpFunctionEnd
+define void @main() #1 {
+entry:
+ %0 = alloca %B, align 4
+ %1 = getelementptr %B, ptr %0, i32 0, i32 2
+ %2 = getelementptr %A, ptr %1, i32 0, i32 1
+ ret void
+}
+
+attributes #1 = { "hlsl.numthreads"="4,8,16" "hlsl.shader"="compute" }
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse-subbyte.ll b/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse-subbyte.ll
index 481bad9a26b7b..280f586891717 100644
--- a/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse-subbyte.ll
+++ b/llvm/test/CodeGen/SPIRV/transcoding/OpBitReverse-subbyte.ll
@@ -19,10 +19,15 @@
; TODO: Add a check to ensure that there's no behavior change of bitreverse operation
; between the LLVM-IR and SPIR-V for i2 and i4
+ at G_res2 = global i2 0
+ at G_res4 = global i4 0
+
define spir_func void @foo(i2 %a, i4 %b) {
entry:
%res2 = tail call i2 @llvm.bitreverse.i2(i2 %a)
+ store i2 %res2, ptr @G_res2
%res4 = tail call i4 @llvm.bitreverse.i4(i4 %b)
+ store i4 %res4, ptr @G_res4
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/OpGenericCastToPtr.ll b/llvm/test/CodeGen/SPIRV/transcoding/OpGenericCastToPtr.ll
index 119dbe14446c1..68f33510b6a8d 100644
--- a/llvm/test/CodeGen/SPIRV/transcoding/OpGenericCastToPtr.ll
+++ b/llvm/test/CodeGen/SPIRV/transcoding/OpGenericCastToPtr.ll
@@ -45,6 +45,12 @@ entry:
%GE = call spir_func ptr addrspace(1) @_Z41__spirv_GenericCastToPtrExplicit_ToGlobalPvi(ptr addrspace(4) %var1, i32 5)
%LE = call spir_func ptr addrspace(3) @_Z40__spirv_GenericCastToPtrExplicit_ToLocalPvi(ptr addrspace(4) %var2, i32 4)
%PE = call spir_func ptr @_Z42__spirv_GenericCastToPtrExplicit_ToPrivatePvi(ptr addrspace(4) %var3, i32 7)
+ store i32 0, ptr addrspace(1) %G, align 4
+ store i8 0, ptr addrspace(3) %L, align 1
+ store i32 0, ptr %P, align 4
+ store i32 0, ptr addrspace(1) %GE, align 4
+ store i8 0, ptr addrspace(3) %LE, align 1
+ store i32 0, ptr %PE, align 4
ret void
}
@@ -70,6 +76,9 @@ entry:
%G = call spir_func ptr addrspace(1) @_Z9to_globalPv(ptr addrspace(4) %var1)
%L = call spir_func ptr addrspace(3) @_Z8to_localPv(ptr addrspace(4) %var2)
%P = call spir_func ptr @_Z10to_privatePv(ptr addrspace(4) %var3)
+ store i32 0, ptr addrspace(1) %G, align 4
+ store i8 0, ptr addrspace(3) %L, align 1
+ store i32 0, ptr %P, align 4
ret void
}
@@ -114,6 +123,12 @@ entry:
%GE = call spir_func ptr addrspace(1) @__spirv_GenericCastToPtrExplicit_ToGlobal(ptr addrspace(4) %var1, i32 5)
%LE = call spir_func ptr addrspace(3) @__spirv_GenericCastToPtrExplicit_ToLocal(ptr addrspace(4) %var2, i32 4)
%PE = call spir_func ptr @__spirv_GenericCastToPtrExplicit_ToPrivate(ptr addrspace(4) %var3, i32 7)
+ store i32 0, ptr addrspace(1) %G, align 4
+ store i8 0, ptr addrspace(3) %L, align 1
+ store i32 0, ptr %P, align 4
+ store i32 0, ptr addrspace(1) %GE, align 4
+ store i8 0, ptr addrspace(3) %LE, align 1
+ store i32 0, ptr %PE, align 4
ret void
}
@@ -139,6 +154,9 @@ entry:
%G = call spir_func ptr addrspace(1) @to_global(ptr addrspace(4) %var1)
%L = call spir_func ptr addrspace(3) @to_local(ptr addrspace(4) %var2)
%P = call spir_func ptr @to_private(ptr addrspace(4) %var3)
+ store i32 0, ptr addrspace(1) %G, align 4
+ store i8 0, ptr addrspace(3) %L, align 1
+ store i32 0, ptr %P, align 4
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/OpPtrCastToGeneric.ll b/llvm/test/CodeGen/SPIRV/transcoding/OpPtrCastToGeneric.ll
index 818243ab19e41..9f08a65c16866 100644
--- a/llvm/test/CodeGen/SPIRV/transcoding/OpPtrCastToGeneric.ll
+++ b/llvm/test/CodeGen/SPIRV/transcoding/OpPtrCastToGeneric.ll
@@ -16,9 +16,13 @@
; CHECK-SPIRV: OpGenericCastToPtr %[[#LocalCharPtr]] %[[#Ptr2]]
; CHECK-SPIRV: OpFunctionEnd
+ at G_p = global ptr addrspace(3) null
+ at G_p2 = global ptr addrspace(3) null
+
define spir_kernel void @foo(ptr addrspace(1) %arg) {
entry:
%p = addrspacecast ptr addrspace(1) %arg to ptr addrspace(3)
+ store ptr addrspace(3) %p, ptr @G_p
ret void
}
@@ -26,5 +30,6 @@ define spir_kernel void @bar(ptr addrspace(1) %arg) {
entry:
%p1 = addrspacecast ptr addrspace(1) %arg to ptr addrspace(4)
%p2 = addrspacecast ptr addrspace(4) %p1 to ptr addrspace(3)
+ store ptr addrspace(3) %p2, ptr @G_p2
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/fcmp.ll b/llvm/test/CodeGen/SPIRV/transcoding/fcmp.ll
index 46eaba9d5ceb1..c752e278927a9 100644
--- a/llvm/test/CodeGen/SPIRV/transcoding/fcmp.ll
+++ b/llvm/test/CodeGen/SPIRV/transcoding/fcmp.ll
@@ -184,6 +184,8 @@
; CHECK-SPIRV: %[[#r89]] = OpUnordered %[[#bool]]
; CHECK-SPIRV: %[[#r90]] = OpUnordered %[[#bool]]
+ at G = global [90 x i1] zeroinitializer
+
define spir_kernel void @testFCmp(float %a, float %b) local_unnamed_addr {
entry:
%r1 = fcmp oeq float %a, %b
@@ -276,5 +278,185 @@ entry:
%r88 = fcmp uno float %a, %b
%r89 = fcmp ninf uno float %a, %b
%r90 = fcmp nsz uno float %a, %b
+ %p1 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 0
+ store i1 %r1, ptr %p1
+ %p2 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 1
+ store i1 %r2, ptr %p2
+ %p3 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 2
+ store i1 %r3, ptr %p3
+ %p4 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 3
+ store i1 %r4, ptr %p4
+ %p5 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 4
+ store i1 %r5, ptr %p5
+ %p6 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 5
+ store i1 %r6, ptr %p6
+ %p7 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 6
+ store i1 %r7, ptr %p7
+ %p8 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 7
+ store i1 %r8, ptr %p8
+ %p9 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 8
+ store i1 %r9, ptr %p9
+ %p10 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 9
+ store i1 %r10, ptr %p10
+ %p11 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 10
+ store i1 %r11, ptr %p11
+ %p12 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 11
+ store i1 %r12, ptr %p12
+ %p13 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 12
+ store i1 %r13, ptr %p13
+ %p14 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 13
+ store i1 %r14, ptr %p14
+ %p15 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 14
+ store i1 %r15, ptr %p15
+ %p16 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 15
+ store i1 %r16, ptr %p16
+ %p17 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 16
+ store i1 %r17, ptr %p17
+ %p18 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 17
+ store i1 %r18, ptr %p18
+ %p19 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 18
+ store i1 %r19, ptr %p19
+ %p20 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 19
+ store i1 %r20, ptr %p20
+ %p21 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 20
+ store i1 %r21, ptr %p21
+ %p22 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 21
+ store i1 %r22, ptr %p22
+ %p23 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 22
+ store i1 %r23, ptr %p23
+ %p24 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 23
+ store i1 %r24, ptr %p24
+ %p25 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 24
+ store i1 %r25, ptr %p25
+ %p26 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 25
+ store i1 %r26, ptr %p26
+ %p27 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 26
+ store i1 %r27, ptr %p27
+ %p28 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 27
+ store i1 %r28, ptr %p28
+ %p29 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 28
+ store i1 %r29, ptr %p29
+ %p30 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 29
+ store i1 %r30, ptr %p30
+ %p31 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 30
+ store i1 %r31, ptr %p31
+ %p32 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 31
+ store i1 %r32, ptr %p32
+ %p33 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 32
+ store i1 %r33, ptr %p33
+ %p34 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 33
+ store i1 %r34, ptr %p34
+ %p35 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 34
+ store i1 %r35, ptr %p35
+ %p36 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 35
+ store i1 %r36, ptr %p36
+ %p37 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 36
+ store i1 %r37, ptr %p37
+ %p38 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 37
+ store i1 %r38, ptr %p38
+ %p39 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 38
+ store i1 %r39, ptr %p39
+ %p40 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 39
+ store i1 %r40, ptr %p40
+ %p41 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 40
+ store i1 %r41, ptr %p41
+ %p42 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 41
+ store i1 %r42, ptr %p42
+ %p43 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 42
+ store i1 %r43, ptr %p43
+ %p44 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 43
+ store i1 %r44, ptr %p44
+ %p45 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 44
+ store i1 %r45, ptr %p45
+ %p46 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 45
+ store i1 %r46, ptr %p46
+ %p47 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 46
+ store i1 %r47, ptr %p47
+ %p48 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 47
+ store i1 %r48, ptr %p48
+ %p49 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 48
+ store i1 %r49, ptr %p49
+ %p50 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 49
+ store i1 %r50, ptr %p50
+ %p51 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 50
+ store i1 %r51, ptr %p51
+ %p52 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 51
+ store i1 %r52, ptr %p52
+ %p53 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 52
+ store i1 %r53, ptr %p53
+ %p54 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 53
+ store i1 %r54, ptr %p54
+ %p55 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 54
+ store i1 %r55, ptr %p55
+ %p56 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 55
+ store i1 %r56, ptr %p56
+ %p57 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 56
+ store i1 %r57, ptr %p57
+ %p58 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 57
+ store i1 %r58, ptr %p58
+ %p59 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 58
+ store i1 %r59, ptr %p59
+ %p60 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 59
+ store i1 %r60, ptr %p60
+ %p61 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 60
+ store i1 %r61, ptr %p61
+ %p62 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 61
+ store i1 %r62, ptr %p62
+ %p63 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 62
+ store i1 %r63, ptr %p63
+ %p64 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 63
+ store i1 %r64, ptr %p64
+ %p65 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 64
+ store i1 %r65, ptr %p65
+ %p66 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 65
+ store i1 %r66, ptr %p66
+ %p67 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 66
+ store i1 %r67, ptr %p67
+ %p68 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 67
+ store i1 %r68, ptr %p68
+ %p69 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 68
+ store i1 %r69, ptr %p69
+ %p70 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 69
+ store i1 %r70, ptr %p70
+ %p71 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 70
+ store i1 %r71, ptr %p71
+ %p72 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 71
+ store i1 %r72, ptr %p72
+ %p73 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 72
+ store i1 %r73, ptr %p73
+ %p74 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 73
+ store i1 %r74, ptr %p74
+ %p75 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 74
+ store i1 %r75, ptr %p75
+ %p76 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 75
+ store i1 %r76, ptr %p76
+ %p77 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 76
+ store i1 %r77, ptr %p77
+ %p78 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 77
+ store i1 %r78, ptr %p78
+ %p79 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 78
+ store i1 %r79, ptr %p79
+ %p80 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 79
+ store i1 %r80, ptr %p80
+ %p81 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 80
+ store i1 %r81, ptr %p81
+ %p82 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 81
+ store i1 %r82, ptr %p82
+ %p83 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 82
+ store i1 %r83, ptr %p83
+ %p84 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 83
+ store i1 %r84, ptr %p84
+ %p85 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 84
+ store i1 %r85, ptr %p85
+ %p86 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 85
+ store i1 %r86, ptr %p86
+ %p87 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 86
+ store i1 %r87, ptr %p87
+ %p88 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 87
+ store i1 %r88, ptr %p88
+ %p89 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 88
+ store i1 %r89, ptr %p89
+ %p90 = getelementptr inbounds [90 x i1], ptr @G, i32 0, i32 89
+ store i1 %r90, ptr %p90
ret void
}
diff --git a/llvm/test/CodeGen/SPIRV/transcoding/spirv-event-null.ll b/llvm/test/CodeGen/SPIRV/transcoding/spirv-event-null.ll
index c8691c32710ad..7658362773218 100644
--- a/llvm/test/CodeGen/SPIRV/transcoding/spirv-event-null.ll
+++ b/llvm/test/CodeGen/SPIRV/transcoding/spirv-event-null.ll
@@ -31,9 +31,12 @@
%StructEvent = type { target("spirv.Event") }
+ at G_r = global target("spirv.Event") poison
+
define spir_kernel void @test_half(ptr addrspace(3) %_arg1, ptr addrspace(1) %_arg2) {
entry:
%r = tail call spir_func target("spirv.Event") @_Z22__spirv_GroupAsyncCopyjPU3AS3Dv2_DF16_PU3AS1KS_mm9ocl_event(i32 2, ptr addrspace(3) %_arg1, ptr addrspace(1) %_arg2, i64 16, i64 10, target("spirv.Event") zeroinitializer)
+ store target("spirv.Event") %r, ptr @G_r
ret void
}
@@ -42,7 +45,6 @@ declare dso_local spir_func target("spirv.Event") @_Z22__spirv_GroupAsyncCopyjPU
; CHECK: OpFunction
; CHECK: OpFunctionParameter
; CHECK: %[[#Src:]] = OpFunctionParameter
-; CHECK: OpVariable %[[#TyStructPtr]] Function
; CHECK: %[[#EventVar:]] = OpVariable %[[#TyEventPtr]] Function
; CHECK: %[[#Dest:]] = OpInBoundsPtrAccessChain
; CHECK: %[[#CopyRes:]] = OpGroupAsyncCopy %[[#TyEvent]] %[[#]] %[[#Dest]] %[[#Src]] %[[#]] %[[#]] %[[#ConstEvent]]
diff --git a/llvm/test/CodeGen/SPIRV/uitofp-with-bool.ll b/llvm/test/CodeGen/SPIRV/uitofp-with-bool.ll
index 46668645f418b..9c8b4070d834d 100644
--- a/llvm/test/CodeGen/SPIRV/uitofp-with-bool.ll
+++ b/llvm/test/CodeGen/SPIRV/uitofp-with-bool.ll
@@ -68,6 +68,27 @@
; SPV-DAG: %[[#ones_64:]] = OpConstantComposite %[[#vec_64]] %[[#one_64]] %[[#one_64]]
; SPV-DAG: %[[#pointer:]] = OpTypePointer CrossWorkgroup %[[#float]]
+ at G_s1 = global i8 0
+ at G_s2 = global i16 0
+ at G_s3 = global i32 0
+ at G_s4 = global i64 0
+ at G_s5 = global <2 x i8> zeroinitializer
+ at G_s6 = global <2 x i16> zeroinitializer
+ at G_s7 = global <2 x i32> zeroinitializer
+ at G_s8 = global <2 x i64> zeroinitializer
+ at G_z1 = global i8 0
+ at G_z2 = global i16 0
+ at G_z3 = global i32 0
+ at G_z4 = global i64 0
+ at G_z5 = global <2 x i8> zeroinitializer
+ at G_z6 = global <2 x i16> zeroinitializer
+ at G_z7 = global <2 x i32> zeroinitializer
+ at G_z8 = global <2 x i64> zeroinitializer
+ at G_ufp1 = global float 0.0
+ at G_ufp2 = global <2 x float> zeroinitializer
+ at G_sfp1 = global float 0.0
+ at G_sfp2 = global <2 x float> zeroinitializer
+
; SPV-DAG: OpFunction
; SPV-DAG: %[[#A:]] = OpFunctionParameter %[[#pointer]]
; SPV-DAG: %[[#B:]] = OpFunctionParameter %[[#]]
@@ -87,47 +108,67 @@ entry:
; SPV-DAG: %[[#s1]] = OpSelect %[[#int_8]] %[[#i1s]] %[[#mone_8]] %[[#zero_8]]
%s1 = sext i1 %i1s to i8
+ store i8 %s1, ptr @G_s1
; SPV-DAG: %[[#s2]] = OpSelect %[[#int_16]] %[[#i1s]] %[[#mone_16]] %[[#zero_16]]
%s2 = sext i1 %i1s to i16
+ store i16 %s2, ptr @G_s2
; SPV-DAG: %[[#s3]] = OpSelect %[[#int_32]] %[[#i1s]] %[[#mone_32]] %[[#zero_32]]
%s3 = sext i1 %i1s to i32
+ store i32 %s3, ptr @G_s3
; SPV-DAG: %[[#s4]] = OpSelect %[[#int_64]] %[[#i1s]] %[[#mone_64]] %[[#zero_64]]
%s4 = sext i1 %i1s to i64
+ store i64 %s4, ptr @G_s4
; SPV-DAG: %[[#s5]] = OpSelect %[[#vec_8]] %[[#i1v]] %[[#mones_8]] %[[#zeros_8]]
%s5 = sext <2 x i1> %i1v to <2 x i8>
+ store <2 x i8> %s5, ptr @G_s5
; SPV-DAG: %[[#s6]] = OpSelect %[[#vec_16]] %[[#i1v]] %[[#mones_16]] %[[#zeros_16]]
%s6 = sext <2 x i1> %i1v to <2 x i16>
+ store <2 x i16> %s6, ptr @G_s6
; SPV-DAG: %[[#s7]] = OpSelect %[[#vec_32]] %[[#i1v]] %[[#mones_32]] %[[#zeros_32]]
%s7 = sext <2 x i1> %i1v to <2 x i32>
+ store <2 x i32> %s7, ptr @G_s7
; SPV-DAG: %[[#s8]] = OpSelect %[[#vec_64]] %[[#i1v]] %[[#mones_64]] %[[#zeros_64]]
%s8 = sext <2 x i1> %i1v to <2 x i64>
+ store <2 x i64> %s8, ptr @G_s8
; SPV-DAG: %[[#z1]] = OpSelect %[[#int_8]] %[[#i1s]] %[[#one_8]] %[[#zero_8]]
%z1 = zext i1 %i1s to i8
+ store i8 %z1, ptr @G_z1
; SPV-DAG: %[[#z2]] = OpSelect %[[#int_16]] %[[#i1s]] %[[#one_16]] %[[#zero_16]]
%z2 = zext i1 %i1s to i16
+ store i16 %z2, ptr @G_z2
; SPV-DAG: %[[#z3]] = OpSelect %[[#int_32]] %[[#i1s]] %[[#one_32]] %[[#zero_32]]
%z3 = zext i1 %i1s to i32
+ store i32 %z3, ptr @G_z3
; SPV-DAG: %[[#z4]] = OpSelect %[[#int_64]] %[[#i1s]] %[[#one_64]] %[[#zero_64]]
%z4 = zext i1 %i1s to i64
+ store i64 %z4, ptr @G_z4
; SPV-DAG: %[[#z5]] = OpSelect %[[#vec_8]] %[[#i1v]] %[[#ones_8]] %[[#zeros_8]]
%z5 = zext <2 x i1> %i1v to <2 x i8>
+ store <2 x i8> %z5, ptr @G_z5
; SPV-DAG: %[[#z6]] = OpSelect %[[#vec_16]] %[[#i1v]] %[[#ones_16]] %[[#zeros_16]]
%z6 = zext <2 x i1> %i1v to <2 x i16>
+ store <2 x i16> %z6, ptr @G_z6
; SPV-DAG: %[[#z7]] = OpSelect %[[#vec_32]] %[[#i1v]] %[[#ones_32]] %[[#zeros_32]]
%z7 = zext <2 x i1> %i1v to <2 x i32>
+ store <2 x i32> %z7, ptr @G_z7
; SPV-DAG: %[[#z8]] = OpSelect %[[#vec_64]] %[[#i1v]] %[[#ones_64]] %[[#zeros_64]]
%z8 = zext <2 x i1> %i1v to <2 x i64>
+ store <2 x i64> %z8, ptr @G_z8
; SPV-DAG: %[[#ufp1_res:]] = OpSelect %[[#int_32]] %[[#i1s]] %[[#one_32]] %[[#zero_32]]
; SPV-DAG: %[[#ufp1]] = OpConvertUToF %[[#float]] %[[#ufp1_res]]
%ufp1 = uitofp i1 %i1s to float
+ store float %ufp1, ptr @G_ufp1
; SPV-DAG: %[[#ufp2_res:]] = OpSelect %[[#vec_32]] %[[#i1v]] %[[#ones_32]] %[[#zeros_32]]
; SPV-DAG: %[[#ufp2]] = OpConvertUToF %[[#vec_float]] %[[#ufp2_res]]
%ufp2 = uitofp <2 x i1> %i1v to <2 x float>
+ store <2 x float> %ufp2, ptr @G_ufp2
; SPV-DAG: %[[#sfp1_res:]] = OpSelect %[[#int_32]] %[[#i1s]] %[[#one_32]] %[[#zero_32]]
; SPV-DAG: %[[#sfp1]] = OpConvertSToF %[[#float]] %[[#sfp1_res]]
%sfp1 = sitofp i1 %i1s to float
+ store float %sfp1, ptr @G_sfp1
; SPV-DAG: %[[#sfp2_res:]] = OpSelect %[[#vec_32]] %[[#i1v]] %[[#ones_32]] %[[#zeros_32]]
; SPV-DAG: %[[#sfp2]] = OpConvertSToF %[[#vec_float]] %[[#sfp2_res]]
%sfp2 = sitofp <2 x i1> %i1v to <2 x float>
+ store <2 x float> %sfp2, ptr @G_sfp2
ret void
}
>From 01f9f4d88226b43f0066e3552a4e2fb71686ba88 Mon Sep 17 00:00:00 2001
From: Steven Perron <stevenperron at google.com>
Date: Wed, 19 Nov 2025 11:51:33 -0500
Subject: [PATCH 2/4] Changes from code review.
---
llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp | 12 +++++++-----
llvm/test/CodeGen/SPIRV/freeze.ll | 2 +-
2 files changed, 8 insertions(+), 6 deletions(-)
diff --git a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
index fd473a45080eb..14b7b9308b351 100644
--- a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
@@ -509,6 +509,9 @@ static bool isConstReg(MachineRegisterInfo *MRI, Register OpReg) {
return false;
}
+// TODO(168736): We should make this either a flag in tabelgen
+// or reduce our dependence on the global registery, so we can remove this
+// function. It can easily be missed when new intrinsics are added.
static bool intrinsicHasSideEffects(Intrinsic::ID ID) {
switch (ID) {
// Intrinsics that do not have side effects.
@@ -572,10 +575,10 @@ static bool intrinsicHasSideEffects(Intrinsic::ID ID) {
}
}
+// TODO(168736): We should make this either a flag in tabelgen
+// or reduce our dependence on the global registery, so we can remove this
+// function. It can easily be missed when new intrinsics are added.
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:
@@ -612,9 +615,8 @@ static bool isOpcodeWithNoSideEffects(unsigned Opcode) {
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) {
+ if (MI.getNumDefs() == 0)
return false;
- }
for (const auto &MO : MI.all_defs()) {
Register Reg = MO.getReg();
diff --git a/llvm/test/CodeGen/SPIRV/freeze.ll b/llvm/test/CodeGen/SPIRV/freeze.ll
index e8c9e933904cc..4f7e7794ed03b 100644
--- a/llvm/test/CodeGen/SPIRV/freeze.ll
+++ b/llvm/test/CodeGen/SPIRV/freeze.ll
@@ -51,4 +51,4 @@ entry:
%val2 = freeze i32 100
%val3 = freeze i32 %val2
ret i32 %val3
-}
\ No newline at end of file
+}
>From 8fb4a310a42a6d1dfdfe0396e5b67a992fdafc88 Mon Sep 17 00:00:00 2001
From: Steven Perron <stevenperron at google.com>
Date: Thu, 20 Nov 2025 09:45:39 -0500
Subject: [PATCH 3/4] Apply suggestions from code review
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit
Co-authored-by: Nathan Gauër <github at keenuts.net>
---
llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp | 6 +++---
1 file changed, 3 insertions(+), 3 deletions(-)
diff --git a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
index 14b7b9308b351..d6e371c178392 100644
--- a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
@@ -510,7 +510,7 @@ static bool isConstReg(MachineRegisterInfo *MRI, Register OpReg) {
}
// TODO(168736): We should make this either a flag in tabelgen
-// 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
// function. It can easily be missed when new intrinsics are added.
static bool intrinsicHasSideEffects(Intrinsic::ID ID) {
switch (ID) {
@@ -576,7 +576,7 @@ static bool intrinsicHasSideEffects(Intrinsic::ID ID) {
}
// TODO(168736): We should make this either a flag in tabelgen
-// 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
// function. It can easily be missed when new intrinsics are added.
static bool isOpcodeWithNoSideEffects(unsigned Opcode) {
switch (Opcode) {
@@ -675,7 +675,7 @@ bool isDead(const MachineInstr &MI, const MachineRegisterInfo &MRI) {
}
void SPIRVInstructionSelector::removeOpNamesForDeadMI(MachineInstr &MI) const {
- // Delete the OpName that uses the result of there is one.
+ // Delete the OpName that uses the result if there is one.
for (const auto &MO : MI.all_defs()) {
Register Reg = MO.getReg();
if (Reg.isPhysical())
>From 094f3231e98c9da838681cedc606ca63ac7c1688 Mon Sep 17 00:00:00 2001
From: Steven Perron <stevenperron at google.com>
Date: Thu, 20 Nov 2025 12:17:21 -0500
Subject: [PATCH 4/4] Fixes based on code review.
---
llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp | 14 ++++++++++----
llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp | 9 +++++++--
2 files changed, 17 insertions(+), 6 deletions(-)
diff --git a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp b/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
index 36facbd05aa67..cd92f4cb5d211 100644
--- a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp
@@ -225,9 +225,15 @@ SPIRVType *SPIRVGlobalRegistry::getOpTypeVoid(MachineIRBuilder &MIRBuilder) {
void SPIRVGlobalRegistry::invalidateMachineInstr(MachineInstr *MI) {
// 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.
+ // - VRegToTypeMap: We cannot remove the definitions of `MI` from
+ // VRegToTypeMap because some calls to invalidateMachineInstr are replacing MI
+ // with another instruction defining the same register. We expect that if MI
+ // is a type instruction, and it is still referenced in VRegToTypeMap, then
+ // those registers are dead or the VRegToTypeMap is out-of-date. We do not
+ // expect passes to ask for the SPIR-V type of a dead register. If the
+ // VRegToTypeMap is out-of-date already, then there was an error before. We
+ // cannot add an assert to verify this because the VRegToTypeMap can be
+ // out-of-date.
// - 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
@@ -337,7 +343,7 @@ Register SPIRVGlobalRegistry::createConstFP(const ConstantFP *CF,
LLT LLTy = LLT::scalar(BitWidth);
Register Res = CurMF->getRegInfo().createGenericVirtualRegister(LLTy);
CurMF->getRegInfo().setRegClass(Res, &SPIRV::fIDRegClass);
- assignFloatTypeToVReg(BitWidth, Res, I, TII);
+ assignSPIRVTypeToVReg(SpvType, Res, *CurMF);
MachineInstr *DepMI = const_cast<MachineInstr *>(SpvType);
MachineIRBuilder MIRBuilder(*DepMI->getParent(), DepMI->getIterator());
diff --git a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
index d6e371c178392..468dbf71235e0 100644
--- a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
@@ -94,7 +94,6 @@ 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;
@@ -512,9 +511,13 @@ static bool isConstReg(MachineRegisterInfo *MRI, Register OpReg) {
// TODO(168736): We should make this either a flag in tabelgen
// or reduce our dependence on the global registry, so we can remove this
// function. It can easily be missed when new intrinsics are added.
+
+// Most SPIR-V instrinsics are considered to have side-effects in their tablegen
+// definition because they are referenced in the global registry. This is a list
+// of intrinsics that have no side effects other than their references in the
+// global registry.
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:
@@ -644,6 +647,8 @@ bool isDead(const MachineInstr &MI, const MachineRegisterInfo &MRI) {
return true;
}
+ // It is possible that the only side effect is that the instruction is referenced in the
+ // global registry. If that is the only side effect, the intrinsic is dead.
if (MI.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
MI.getOpcode() == TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS) {
const auto &Intr = cast<GIntrinsic>(MI);
More information about the llvm-commits
mailing list