[llvm] [AMDGPU] Refresh RegisterClassInfo after reserving WWM registers (PR #221849)
Michael Halkenhäuser via llvm-commits
llvm-commits at lists.llvm.org
Tue Sep 8 04:46:35 PDT 2026
https://github.com/mhalk updated https://github.com/llvm/llvm-project/pull/221849
>From 3c00823ca37ba615e2988386cba414a219a6bbef Mon Sep 17 00:00:00 2001
From: Michael Halkenhaeuser <MichaelGerald.Halkenhauser at amd.com>
Date: Mon, 7 Sep 2026 23:02:58 +0000
Subject: [PATCH 1/3] [AMDGPU] Refresh RegisterClassInfo after reserving WWM
registers
AMDGPUReserveWWMRegs reserves the VGPRs the WWM allocator picked for SGPR
spill lanes and clears the per-lane VGPR mask, but updates neither
MachineRegisterInfo nor the shared RegisterClassInfo. RegAllocBase
re-freezes the reserved set, but owns its RegisterClassInfo by value, so
only its private copy is refreshed, and a later user of the analysis is
offered an allocation order that still lists the reserved register.
AMDGPURewriteAGPRCopyMFMA became such a user in #216510: its spill slot
elimination assigned that register, which asserts in VirtRegMap and, in a
build without assertions, silently clobbers the spilled SGPRs. #216510 was
reverted in #221749; this is a prerequisite for relanding it.
Update MRI and the analysis at the end of the pass, as SIPreAllocateWWMRegs
already does, after clearPerLaneVGPRAllocMask() since the mask is itself an
input to getReservedRegs().
AI-assisted.
---
.../Target/AMDGPU/AMDGPUReserveWWMRegs.cpp | 33 ++++++++++++++-----
1 file changed, 25 insertions(+), 8 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUReserveWWMRegs.cpp b/llvm/lib/Target/AMDGPU/AMDGPUReserveWWMRegs.cpp
index cc3df121635a2..54f637e589c87 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUReserveWWMRegs.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUReserveWWMRegs.cpp
@@ -18,7 +18,9 @@
#include "AMDGPU.h"
#include "SIMachineFunctionInfo.h"
#include "llvm/CodeGen/MachineFunctionPass.h"
+#include "llvm/CodeGen/RegisterClassInfo.h"
#include "llvm/CodeGen/VirtRegMap.h"
+#include "llvm/InitializePasses.h"
using namespace llvm;
@@ -39,35 +41,44 @@ class AMDGPUReserveWWMRegsLegacy : public MachineFunctionPass {
}
void getAnalysisUsage(AnalysisUsage &AU) const override {
+ AU.addRequired<MachineRegisterClassInfoWrapperPass>();
AU.setPreservesAll();
MachineFunctionPass::getAnalysisUsage(AU);
}
};
class AMDGPUReserveWWMRegs {
+ RegisterClassInfo &RCI;
+
public:
+ explicit AMDGPUReserveWWMRegs(RegisterClassInfo &RCI) : RCI(RCI) {}
+
bool run(MachineFunction &MF);
};
} // End anonymous namespace.
-INITIALIZE_PASS(AMDGPUReserveWWMRegsLegacy, DEBUG_TYPE,
- "AMDGPU Reserve WWM Registers", false, false)
+INITIALIZE_PASS_BEGIN(AMDGPUReserveWWMRegsLegacy, DEBUG_TYPE,
+ "AMDGPU Reserve WWM Registers", false, false)
+INITIALIZE_PASS_DEPENDENCY(MachineRegisterClassInfoWrapperPass)
+INITIALIZE_PASS_END(AMDGPUReserveWWMRegsLegacy, DEBUG_TYPE,
+ "AMDGPU Reserve WWM Registers", false, false)
char AMDGPUReserveWWMRegsLegacy::ID = 0;
char &llvm::AMDGPUReserveWWMRegsLegacyID = AMDGPUReserveWWMRegsLegacy::ID;
bool AMDGPUReserveWWMRegsLegacy::runOnMachineFunction(MachineFunction &MF) {
- return AMDGPUReserveWWMRegs().run(MF);
+ auto &RCI = getAnalysis<MachineRegisterClassInfoWrapperPass>().getRCI();
+ return AMDGPUReserveWWMRegs(RCI).run(MF);
}
PreservedAnalyses
AMDGPUReserveWWMRegsPass::run(MachineFunction &MF,
- MachineFunctionAnalysisManager &) {
- AMDGPUReserveWWMRegs().run(MF);
- // TODO: This should abandon RegisterClassInfo once it is turned into an
- // analysis.
+ MachineFunctionAnalysisManager &MFAM) {
+ auto &RCI = MFAM.getResult<MachineRegisterClassAnalysis>(MF);
+ AMDGPUReserveWWMRegs(RCI).run(MF);
+ // RegisterClassInfo was updated in place, so it need not be abandoned.
return PreservedAnalyses::all();
}
@@ -97,7 +108,7 @@ bool AMDGPUReserveWWMRegs::run(MachineFunction &MF) {
// The renamable flag can't be set for reserved registers. Reset the flag for
// MOs involving wwm-regs as they will be reserved during vgpr-regalloc
// pipeline.
- const MachineRegisterInfo &MRI = MF.getRegInfo();
+ MachineRegisterInfo &MRI = MF.getRegInfo();
for (Register Reg : MFI->getWWMReservedRegs()) {
for (MachineOperand &MO : MRI.reg_operands(Reg))
MO.setIsRenamable(false);
@@ -106,5 +117,11 @@ bool AMDGPUReserveWWMRegs::run(MachineFunction &MF) {
// Now clear the PerLaneVGPRMask earlier set during wwm-regalloc.
MFI->clearPerLaneVGPRAllocMask();
+ // Reserving WWM registers and clearing the per-lane mask both change
+ // getReservedRegs(). Update MRI and the shared RegisterClassInfo; the
+ // following register allocator refreshes only its private instance.
+ MRI.freezeReservedRegs();
+ RCI.updateReservedRegs(MRI.getReservedRegs());
+
return Changed;
}
>From 5867e72d4cae63ba2d809ccdf431edfd216b304f Mon Sep 17 00:00:00 2001
From: Michael Halkenhaeuser <MichaelGerald.Halkenhauser at amd.com>
Date: Mon, 7 Sep 2026 23:02:58 +0000
Subject: [PATCH 2/3] [AMDGPU] Add test for WWM register reservation vs
RegisterClassInfo
Checks that the register reserved for the WWM SGPR spill lanes is not also
handed to a virtual register by AMDGPURewriteAGPRCopyMFMA. The test only
discriminates once #216510 is relanded; until then it passes either way.
AI-assisted.
---
...ite-vgpr-mfma-to-agpr-reserved-wwm-vgpr.ll | 121 ++++++++++++++++++
1 file changed, 121 insertions(+)
create mode 100644 llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-reserved-wwm-vgpr.ll
diff --git a/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-reserved-wwm-vgpr.ll b/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-reserved-wwm-vgpr.ll
new file mode 100644
index 0000000000000..2bd46639c7c93
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-reserved-wwm-vgpr.ll
@@ -0,0 +1,121 @@
+; RUN: llc -mtriple=amdgpu9.50-amd-amdhsa < %s \
+; RUN: | FileCheck %s --implicit-check-not='v_mov_b32_e32 v23,'
+; RUN: llc -enable-new-pm -mtriple=amdgpu9.50-amd-amdhsa < %s \
+; RUN: | FileCheck %s --implicit-check-not='v_mov_b32_e32 v23,'
+
+; The "amdgpu-num-vgpr"="24" attribute leaves v0..v23 addressable and the WWM
+; allocator takes v23 for the SGPR spill lanes. AMDGPUReserveWWMRegs used to
+; reserve it without refreshing the shared RegisterClassInfo, so
+; AMDGPURewriteAGPRCopyMFMA was still offered v23 and its spill slot elimination
+; assigned it: an assertion in VirtRegMap with assertions on, a silent clobber
+; of the spilled SGPRs without.
+;
+; The implicit-check-not is what catches that clobber without assertions. The
+; MFMA and ScratchSize checks stop the test passing vacuously if the rewrite
+; stops running or the kernel stops spilling.
+
+; CHECK-LABEL: {{^}}reserved_wwm_vgpr_not_in_alloc_order:
+; CHECK: v_writelane_b32 v23,
+; CHECK: v_mfma_f32_32x32x16_bf16 a[
+; CHECK: v_readlane_b32 {{s[0-9]+}}, v23,
+; CHECK: ; NumVgprs: 24
+; CHECK: ; ScratchSize: {{[1-9][0-9]*}}
+declare noundef range(i32 0, 1024) i32 @llvm.amdgcn.workitem.id.x() #0
+declare <4 x i32> @llvm.amdgcn.raw.ptr.buffer.load.v4i32(ptr addrspace(8) readonly captures(none), i32, i32, i32 immarg) #1
+declare void @llvm.amdgcn.sched.barrier(i32 immarg) #2
+declare <16 x float> @llvm.amdgcn.mfma.f32.32x32x16.bf16(<8 x bfloat>, <8 x bfloat>, <16 x float>, i32 immarg, i32 immarg, i32 immarg) #3
+
+define amdgpu_kernel void @reserved_wwm_vgpr_not_in_alloc_order(i32 %arg6, i32 %arg11, i1 %arg15, i1 %arg16, i1 %arg21, ptr addrspace(3) %arg23, i1 %arg26, i1 %arg27, ptr addrspace(3) %arg31, ptr addrspace(3) %arg35, i32 %arg36, i1 %arg37, i1 %arg38, <4 x bfloat> %arg44, ptr addrspace(3) %arg46, <8 x bfloat> %phi123, <8 x bfloat> %bitcast214, i1 %phi119, <8 x i8> %bitcast252, <8 x bfloat> %sitofp, <8 x bfloat> %bitcast216, <8 x bfloat> %bitcast218, <8 x bfloat> %bitcast212) #4 {
+bbl:
+ br label %bbl97
+
+bbl97: ; preds = %bbl97, %bbl
+ %phi = phi i1 [ false, %bbl ], [ %and171, %bbl97 ]
+ %phi98 = phi i1 [ false, %bbl ], [ %and175, %bbl97 ]
+ %phi99 = phi i1 [ false, %bbl ], [ %and179, %bbl97 ]
+ %phi101 = phi i1 [ false, %bbl ], [ %arg15, %bbl97 ]
+ %phi102 = phi i1 [ false, %bbl ], [ %arg26, %bbl97 ]
+ %phi103 = phi i1 [ false, %bbl ], [ %arg37, %bbl97 ]
+ %phi104 = phi i1 [ false, %bbl ], [ %arg16, %bbl97 ]
+ %phi106 = phi i32 [ 0, %bbl ], [ 1, %bbl97 ]
+ %phi1071 = phi i32 [ 0, %bbl ], [ 1, %bbl97 ]
+ %phi112 = phi <4 x i32> [ zeroinitializer, %bbl ], [ %call184, %bbl97 ]
+ %phi113 = phi <4 x i32> [ zeroinitializer, %bbl ], [ splat (i32 1), %bbl97 ]
+ %phi116 = phi i1 [ false, %bbl ], [ %and250, %bbl97 ]
+ %phi117 = phi i1 [ false, %bbl ], [ %and256, %bbl97 ]
+ %phi118 = phi i1 [ false, %bbl ], [ %arg38, %bbl97 ]
+ %phi121 = phi i32 [ 0, %bbl ], [ 1, %bbl97 ]
+ %phi122 = phi <8 x bfloat> [ zeroinitializer, %bbl ], [ %sitofp272, %bbl97 ]
+ %phi128 = phi <16 x float> [ zeroinitializer, %bbl ], [ %call281, %bbl97 ]
+ %bitcast162 = bitcast <4 x i32> %phi112 to <8 x bfloat>
+ %select163 = select i1 %arg21, <8 x bfloat> %bitcast162, <8 x bfloat> zeroinitializer
+ %select = select i1 %phi, <8 x bfloat> splat (bfloat 1.000000e+00), <8 x bfloat> zeroinitializer
+ %getelementptr164 = getelementptr [2 x i8], ptr addrspace(3) null, i32 %phi1071
+ store <8 x bfloat> %select, ptr addrspace(3) %getelementptr164, align 16
+ %select159 = select i1 %phi98, <8 x bfloat> splat (bfloat 1.000000e+00), <8 x bfloat> zeroinitializer
+ store <8 x bfloat> %select159, ptr addrspace(3) null, align 16
+ %bitcast160 = bitcast <4 x i32> %phi113 to <8 x bfloat>
+ %select161 = select i1 %phi99, <8 x bfloat> %bitcast160, <8 x bfloat> zeroinitializer
+ store <8 x bfloat> %select161, ptr addrspace(3) %arg23, align 16
+ store <8 x bfloat> %select163, ptr addrspace(3) null, align 16
+ %call1 = tail call i32 @llvm.amdgcn.workitem.id.x()
+ %icmp = icmp slt i32 %call1, 1
+ %icmp170 = icmp slt i32 %phi106, 1
+ %and171 = and i1 %icmp, %icmp170
+ %icmp73 = icmp slt i32 %call1, %arg6
+ %and175 = and i1 %icmp73, %icmp170
+ %and = and i32 %call1, 1020
+ %icmp75 = icmp slt i32 %and, 1
+ %and179 = and i1 %icmp75, %icmp170
+ %call184 = tail call <4 x i32> @llvm.amdgcn.raw.ptr.buffer.load.v4i32(ptr addrspace(8) null, i32 0, i32 0, i32 0)
+ %select219 = select i1 %phi104, <8 x bfloat> %bitcast218, <8 x bfloat> zeroinitializer
+ %select213 = select i1 %phi101, <8 x bfloat> %sitofp, <8 x bfloat> zeroinitializer
+ store <8 x bfloat> %select213, ptr addrspace(3) %arg23, align 16
+ %select215 = select i1 %phi102, <8 x bfloat> %bitcast214, <8 x bfloat> zeroinitializer
+ store <8 x bfloat> %select215, ptr addrspace(3) null, align 16
+ %select217 = select i1 %phi103, <8 x bfloat> %bitcast216, <8 x bfloat> zeroinitializer
+ store <8 x bfloat> %select217, ptr addrspace(3) %arg31, align 16
+ store <8 x bfloat> %select219, ptr addrspace(3) %arg35, align 16
+ %select227 = select i1 %phi118, <8 x bfloat> %phi123, <8 x bfloat> zeroinitializer
+ %shufflevector230 = shufflevector <8 x bfloat> %select227, <8 x bfloat> zeroinitializer, <4 x i32> <i32 poison, i32 poison, i32 poison, i32 7>
+ %select225 = select i1 %phi116, <8 x bfloat> %bitcast212, <8 x bfloat> zeroinitializer
+ %select226 = select i1 %phi117, <8 x bfloat> splat (bfloat 1.000000e+00), <8 x bfloat> zeroinitializer
+ %shufflevector229 = shufflevector <8 x bfloat> %select225, <8 x bfloat> %select226, <4 x i32> <i32 7, i32 15, i32 poison, i32 poison>
+ %shufflevector231 = shufflevector <4 x bfloat> %shufflevector229, <4 x bfloat> %shufflevector230, <4 x i32> <i32 0, i32 1, i32 7, i32 poison>
+ %select228 = select i1 %arg27, <8 x bfloat> %phi122, <8 x bfloat> zeroinitializer
+ %shufflevector232 = shufflevector <8 x bfloat> %select228, <8 x bfloat> zeroinitializer, <4 x i32> <i32 poison, i32 poison, i32 poison, i32 7>
+ %shufflevector233 = shufflevector <4 x bfloat> %shufflevector231, <4 x bfloat> %shufflevector232, <4 x i32> <i32 0, i32 1, i32 2, i32 7>
+ store <4 x bfloat> %shufflevector233, ptr addrspace(3) null, align 8
+ %icmp249 = icmp slt i32 %phi121, %arg6
+ %and250 = and i1 %phi119, %icmp249
+ %and64 = and i32 %call1, 1
+ %icmp783 = icmp slt i32 %and64, 1
+ %icmp255 = icmp slt i32 %phi121, %arg36
+ %and256 = and i1 %icmp783, %icmp255
+ %sitofp272 = sitofp <8 x i8> %bitcast252 to <8 x bfloat>
+ %call153 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x16.bf16(<8 x bfloat> splat (bfloat 1.000000e+00), <8 x bfloat> zeroinitializer, <16 x float> %phi128, i32 0, i32 0, i32 0)
+ %call281 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x16.bf16(<8 x bfloat> zeroinitializer, <8 x bfloat> zeroinitializer, <16 x float> %call153, i32 0, i32 0, i32 0)
+ tail call void @llvm.amdgcn.sched.barrier(i32 0)
+ br i1 %arg15, label %bbl97, label %bbl290
+
+bbl290: ; preds = %bbl97
+ %select292 = select i1 %and171, <8 x bfloat> splat (bfloat 1.000000e+00), <8 x bfloat> zeroinitializer
+ store <8 x bfloat> %select292, ptr addrspace(3) null, align 16
+ %select294 = select i1 %and175, <8 x bfloat> splat (bfloat 1.000000e+00), <8 x bfloat> zeroinitializer
+ store <8 x bfloat> %select294, ptr addrspace(3) %arg46, align 16
+ %select296 = select i1 %and179, <8 x bfloat> splat (bfloat 1.000000e+00), <8 x bfloat> zeroinitializer
+ store <8 x bfloat> %select296, ptr addrspace(3) null, align 16
+ %getelementptr301 = getelementptr [2 x i8], ptr addrspace(3) null, i32 %arg11
+ store <8 x bfloat> zeroinitializer, ptr addrspace(3) %getelementptr301, align 16
+ %select314 = select i1 %and256, <8 x bfloat> splat (bfloat 1.000000e+00), <8 x bfloat> zeroinitializer
+ %shufflevector316 = shufflevector <8 x bfloat> zeroinitializer, <8 x bfloat> %select314, <4 x i32> <i32 0, i32 8, i32 poison, i32 poison>
+ %shufflevector318 = shufflevector <4 x bfloat> %shufflevector316, <4 x bfloat> zeroinitializer, <4 x i32> <i32 0, i32 1, i32 4, i32 poison>
+ %shufflevector319 = shufflevector <4 x bfloat> %shufflevector318, <4 x bfloat> zeroinitializer, <4 x i32> <i32 0, i32 1, i32 2, i32 4>
+ ret void
+}
+
+attributes #0 = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
+attributes #1 = { nocallback nofree nosync nounwind willreturn memory(argmem: read) }
+attributes #2 = { convergent nocallback nofree nounwind willreturn }
+attributes #3 = { convergent nocallback nocreateundeforpoison nofree nosync nounwind willreturn memory(none) }
+attributes #4 = { "amdgpu-num-vgpr"="24" }
>From 78438c82dcbd8a3262a58cbeb46ed6a0a38854cf Mon Sep 17 00:00:00 2001
From: Michael Halkenhaeuser <MichaelGerald.Halkenhauser at amd.com>
Date: Tue, 8 Sep 2026 11:45:37 +0000
Subject: [PATCH 3/3] fixup! [AMDGPU] Add test for WWM register reservation vs
RegisterClassInfo
---
...ite-vgpr-mfma-to-agpr-reserved-wwm-vgpr.ll | 446 ++++++++++++++----
1 file changed, 352 insertions(+), 94 deletions(-)
diff --git a/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-reserved-wwm-vgpr.ll b/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-reserved-wwm-vgpr.ll
index 2bd46639c7c93..2f9c62fd283d0 100644
--- a/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-reserved-wwm-vgpr.ll
+++ b/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-reserved-wwm-vgpr.ll
@@ -1,121 +1,379 @@
-; RUN: llc -mtriple=amdgpu9.50-amd-amdhsa < %s \
-; RUN: | FileCheck %s --implicit-check-not='v_mov_b32_e32 v23,'
-; RUN: llc -enable-new-pm -mtriple=amdgpu9.50-amd-amdhsa < %s \
-; RUN: | FileCheck %s --implicit-check-not='v_mov_b32_e32 v23,'
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc -mtriple=amdgpu9.50-amd-amdhsa < %s | FileCheck %s
+; RUN: llc -enable-new-pm -mtriple=amdgpu9.50-amd-amdhsa < %s | FileCheck %s
-; The "amdgpu-num-vgpr"="24" attribute leaves v0..v23 addressable and the WWM
-; allocator takes v23 for the SGPR spill lanes. AMDGPUReserveWWMRegs used to
-; reserve it without refreshing the shared RegisterClassInfo, so
-; AMDGPURewriteAGPRCopyMFMA was still offered v23 and its spill slot elimination
-; assigned it: an assertion in VirtRegMap with assertions on, a silent clobber
-; of the spilled SGPRs without.
-;
-; The implicit-check-not is what catches that clobber without assertions. The
-; MFMA and ScratchSize checks stop the test passing vacuously if the rewrite
-; stops running or the kernel stops spilling.
+; "amdgpu-waves-per-eu"="8,8" leaves v0..v31 addressable and the WWM allocator
+; takes v31 for the SGPR spill lanes. AMDGPUReserveWWMRegs used to reserve it
+; without refreshing the shared RegisterClassInfo, so AMDGPURewriteAGPRCopyMFMA
+; was still offered v31 and its spill slot elimination assigned it: an assertion
+; in VirtRegMap with assertions on, a silent clobber of the spilled SGPRs
+; without.
-; CHECK-LABEL: {{^}}reserved_wwm_vgpr_not_in_alloc_order:
-; CHECK: v_writelane_b32 v23,
-; CHECK: v_mfma_f32_32x32x16_bf16 a[
-; CHECK: v_readlane_b32 {{s[0-9]+}}, v23,
-; CHECK: ; NumVgprs: 24
-; CHECK: ; ScratchSize: {{[1-9][0-9]*}}
-declare noundef range(i32 0, 1024) i32 @llvm.amdgcn.workitem.id.x() #0
-declare <4 x i32> @llvm.amdgcn.raw.ptr.buffer.load.v4i32(ptr addrspace(8) readonly captures(none), i32, i32, i32 immarg) #1
-declare void @llvm.amdgcn.sched.barrier(i32 immarg) #2
-declare <16 x float> @llvm.amdgcn.mfma.f32.32x32x16.bf16(<8 x bfloat>, <8 x bfloat>, <16 x float>, i32 immarg, i32 immarg, i32 immarg) #3
+declare <4 x i32> @llvm.amdgcn.raw.ptr.buffer.load.v4i32(ptr addrspace(8) readonly captures(none), i32, i32, i32 immarg)
+declare <16 x float> @llvm.amdgcn.mfma.f32.32x32x16.bf16(<8 x bfloat>, <8 x bfloat>, <16 x float>, i32 immarg, i32 immarg, i32 immarg)
-define amdgpu_kernel void @reserved_wwm_vgpr_not_in_alloc_order(i32 %arg6, i32 %arg11, i1 %arg15, i1 %arg16, i1 %arg21, ptr addrspace(3) %arg23, i1 %arg26, i1 %arg27, ptr addrspace(3) %arg31, ptr addrspace(3) %arg35, i32 %arg36, i1 %arg37, i1 %arg38, <4 x bfloat> %arg44, ptr addrspace(3) %arg46, <8 x bfloat> %phi123, <8 x bfloat> %bitcast214, i1 %phi119, <8 x i8> %bitcast252, <8 x bfloat> %sitofp, <8 x bfloat> %bitcast216, <8 x bfloat> %bitcast218, <8 x bfloat> %bitcast212) #4 {
+define amdgpu_kernel void @reserved_wwm_vgpr_not_in_alloc_order(i32 %arg6, i1 %arg15, i1 %arg16, i1 %arg21, ptr addrspace(3) %arg23, ptr addrspace(3) %arg33, i1 %arg37, i1 %arg38, <8 x bfloat> %phi123, i1 %and256, <8 x bfloat> %sitofp, i64 %0, i1 %phi103, <8 x bfloat> %bitcast160, <8 x bfloat> %bitcast218, <8 x bfloat> %select294, i1 %icmp) #0 {
+; CHECK-LABEL: reserved_wwm_vgpr_not_in_alloc_order:
+; CHECK: ; %bb.0: ; %bbl
+; CHECK-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x0
+; CHECK-NEXT: s_load_dword s9, s[4:5], 0x30
+; CHECK-NEXT: s_load_dword s24, s[4:5], 0x90
+; CHECK-NEXT: s_load_dword s8, s[4:5], 0x10
+; CHECK-NEXT: s_load_dwordx4 s[12:15], s[4:5], 0x20
+; CHECK-NEXT: ; implicit-def: $vgpr31 : SGPR spill to VGPR lane
+; CHECK-NEXT: v_mov_b32_e32 v30, 0
+; CHECK-NEXT: s_waitcnt lgkmcnt(0)
+; CHECK-NEXT: s_bitcmp1_b32 s1, 0
+; CHECK-NEXT: s_cselect_b64 s[6:7], -1, 0
+; CHECK-NEXT: s_xor_b64 s[16:17], s[6:7], -1
+; CHECK-NEXT: v_writelane_b32 v31, s12, 0
+; CHECK-NEXT: s_bitcmp1_b32 s1, 8
+; CHECK-NEXT: s_cselect_b64 s[10:11], -1, 0
+; CHECK-NEXT: v_writelane_b32 v31, s13, 1
+; CHECK-NEXT: v_writelane_b32 v31, s14, 2
+; CHECK-NEXT: v_writelane_b32 v31, s15, 3
+; CHECK-NEXT: s_bitcmp1_b32 s1, 16
+; CHECK-NEXT: v_writelane_b32 v31, s10, 4
+; CHECK-NEXT: s_cselect_b64 vcc, -1, 0
+; CHECK-NEXT: s_bitcmp1_b32 s8, 0
+; CHECK-NEXT: v_writelane_b32 v31, s11, 5
+; CHECK-NEXT: s_cselect_b64 s[10:11], -1, 0
+; CHECK-NEXT: s_bitcmp1_b32 s8, 8
+; CHECK-NEXT: s_load_dwordx4 s[20:23], s[4:5], 0x40
+; CHECK-NEXT: s_load_dword s1, s[4:5], 0x50
+; CHECK-NEXT: s_load_dword s8, s[4:5], 0x58
+; CHECK-NEXT: s_cselect_b64 s[36:37], -1, 0
+; CHECK-NEXT: s_bitcmp1_b32 s9, 0
+; CHECK-NEXT: s_cselect_b64 s[38:39], -1, 0
+; CHECK-NEXT: v_writelane_b32 v31, s10, 6
+; CHECK-NEXT: s_waitcnt lgkmcnt(0)
+; CHECK-NEXT: s_bitcmp1_b32 s8, 0
+; CHECK-NEXT: s_cselect_b64 s[40:41], -1, 0
+; CHECK-NEXT: s_bitcmp1_b32 s24, 0
+; CHECK-NEXT: s_mov_b32 s24, 0
+; CHECK-NEXT: v_writelane_b32 v31, s11, 7
+; CHECK-NEXT: s_load_dwordx8 s[8:15], s[4:5], 0x60
+; CHECK-NEXT: s_mov_b32 s25, s24
+; CHECK-NEXT: s_mov_b32 s26, s24
+; CHECK-NEXT: s_mov_b32 s27, s24
+; CHECK-NEXT: v_mov_b64_e32 v[0:1], s[24:25]
+; CHECK-NEXT: v_mov_b64_e32 v[2:3], s[26:27]
+; CHECK-NEXT: s_cselect_b64 s[42:43], -1, 0
+; CHECK-NEXT: s_lshl_b32 s60, s1, 1
+; CHECK-NEXT: scratch_store_dwordx4 off, v[0:3], off offset:64 ; 16-byte Folded Spill
+; CHECK-NEXT: v_mov_b32_e32 v4, v30
+; CHECK-NEXT: v_mov_b32_e32 v5, v30
+; CHECK-NEXT: v_mov_b32_e32 v0, 0
+; CHECK-NEXT: v_mov_b32_e32 v1, v30
+; CHECK-NEXT: v_mov_b32_e32 v2, v30
+; CHECK-NEXT: v_mov_b32_e32 v3, v30
+; CHECK-NEXT: v_mov_b32_e32 v6, v30
+; CHECK-NEXT: v_mov_b32_e32 v7, v30
+; CHECK-NEXT: v_mov_b32_e32 v8, v30
+; CHECK-NEXT: v_mov_b32_e32 v9, v30
+; CHECK-NEXT: v_mov_b32_e32 v10, v30
+; CHECK-NEXT: v_mov_b32_e32 v11, v30
+; CHECK-NEXT: v_mov_b32_e32 v12, v30
+; CHECK-NEXT: v_mov_b32_e32 v13, v30
+; CHECK-NEXT: v_mov_b32_e32 v14, v30
+; CHECK-NEXT: v_mov_b32_e32 v15, v30
+; CHECK-NEXT: v_accvgpr_write_b32 a31, v15
+; CHECK-NEXT: v_accvgpr_write_b32 a30, v14
+; CHECK-NEXT: v_accvgpr_write_b32 a29, v13
+; CHECK-NEXT: v_accvgpr_write_b32 a28, v12
+; CHECK-NEXT: v_accvgpr_write_b32 a27, v11
+; CHECK-NEXT: v_accvgpr_write_b32 a26, v10
+; CHECK-NEXT: v_accvgpr_write_b32 a25, v9
+; CHECK-NEXT: v_accvgpr_write_b32 a24, v8
+; CHECK-NEXT: v_accvgpr_write_b32 a23, v7
+; CHECK-NEXT: v_accvgpr_write_b32 a22, v6
+; CHECK-NEXT: v_accvgpr_write_b32 a21, v5
+; CHECK-NEXT: v_accvgpr_write_b32 a20, v4
+; CHECK-NEXT: v_accvgpr_write_b32 a19, v3
+; CHECK-NEXT: v_accvgpr_write_b32 a18, v2
+; CHECK-NEXT: v_accvgpr_write_b32 a17, v1
+; CHECK-NEXT: v_accvgpr_write_b32 a16, v0
+; CHECK-NEXT: v_mov_b32_e32 v0, s60
+; CHECK-NEXT: s_mov_b64 s[52:53], -1
+; CHECK-NEXT: s_mov_b64 s[34:35], s[4:5]
+; CHECK-NEXT: s_mov_b64 s[4:5], s[16:17]
+; CHECK-NEXT: s_mov_b64 s[54:55], 0
+; CHECK-NEXT: s_mov_b64 s[58:59], 0
+; CHECK-NEXT: s_mov_b64 s[56:57], 0
+; CHECK-NEXT: s_mov_b64 s[44:45], 0
+; CHECK-NEXT: s_mov_b64 s[46:47], 0
+; CHECK-NEXT: s_mov_b32 s1, s24
+; CHECK-NEXT: s_mov_b64 s[48:49], 0
+; CHECK-NEXT: s_mov_b64 s[50:51], 0
+; CHECK-NEXT: v_accvgpr_write_b32 a0, 0
+; CHECK-NEXT: v_accvgpr_write_b32 a1, v30
+; CHECK-NEXT: v_accvgpr_write_b32 a2, v30
+; CHECK-NEXT: v_accvgpr_write_b32 a3, v30
+; CHECK-NEXT: v_accvgpr_write_b32 a4, v30
+; CHECK-NEXT: v_accvgpr_write_b32 a5, v30
+; CHECK-NEXT: v_accvgpr_write_b32 a6, v30
+; CHECK-NEXT: v_accvgpr_write_b32 a7, v30
+; CHECK-NEXT: v_accvgpr_write_b32 a8, v30
+; CHECK-NEXT: v_accvgpr_write_b32 a9, v30
+; CHECK-NEXT: v_accvgpr_write_b32 a10, v30
+; CHECK-NEXT: v_accvgpr_write_b32 a11, v30
+; CHECK-NEXT: v_accvgpr_write_b32 a12, v30
+; CHECK-NEXT: v_accvgpr_write_b32 a13, v30
+; CHECK-NEXT: v_accvgpr_write_b32 a14, v30
+; CHECK-NEXT: v_accvgpr_write_b32 a15, v30
+; CHECK-NEXT: s_mov_b32 s33, s24
+; CHECK-NEXT: v_mov_b32_e32 v16, 0
+; CHECK-NEXT: v_mov_b32_e32 v17, v30
+; CHECK-NEXT: v_mov_b32_e32 v18, v30
+; CHECK-NEXT: v_mov_b32_e32 v19, v30
+; CHECK-NEXT: s_mov_b32 s62, s24
+; CHECK-NEXT: s_mov_b32 s63, s24
+; CHECK-NEXT: s_mov_b32 s64, s24
+; CHECK-NEXT: s_mov_b32 s65, s24
+; CHECK-NEXT: scratch_store_dword off, v0, off offset:80 ; 4-byte Folded Spill
+; CHECK-NEXT: v_mov_b32_e32 v0, s2
+; CHECK-NEXT: scratch_store_dword off, v0, off offset:84 ; 4-byte Folded Spill
+; CHECK-NEXT: .p2align 5, , 4
+; CHECK-NEXT: .LBB0_1: ; %bbl97
+; CHECK-NEXT: ; =>This Inner Loop Header: Depth=1
+; CHECK-NEXT: v_accvgpr_read_b32 v0, a0
+; CHECK-NEXT: v_accvgpr_read_b32 v1, a1
+; CHECK-NEXT: v_accvgpr_read_b32 v2, a2
+; CHECK-NEXT: v_accvgpr_read_b32 v3, a3
+; CHECK-NEXT: v_accvgpr_read_b32 v4, a4
+; CHECK-NEXT: v_accvgpr_read_b32 v5, a5
+; CHECK-NEXT: v_accvgpr_read_b32 v6, a6
+; CHECK-NEXT: v_accvgpr_read_b32 v7, a7
+; CHECK-NEXT: v_accvgpr_read_b32 v8, a8
+; CHECK-NEXT: v_accvgpr_read_b32 v9, a9
+; CHECK-NEXT: v_accvgpr_read_b32 v10, a10
+; CHECK-NEXT: v_accvgpr_read_b32 v11, a11
+; CHECK-NEXT: v_accvgpr_read_b32 v12, a12
+; CHECK-NEXT: v_accvgpr_read_b32 v13, a13
+; CHECK-NEXT: v_accvgpr_read_b32 v14, a14
+; CHECK-NEXT: v_accvgpr_read_b32 v15, a15
+; CHECK-NEXT: scratch_load_dwordx4 a[8:11], off, off offset:64 ; 16-byte Folded Reload
+; CHECK-NEXT: scratch_load_dword v20, off, off offset:84 ; 4-byte Folded Reload
+; CHECK-NEXT: s_and_b64 s[54:55], s[54:55], exec
+; CHECK-NEXT: s_cselect_b32 s60, 0x3f803f80, 0
+; CHECK-NEXT: s_and_b64 s[54:55], s[58:59], exec
+; CHECK-NEXT: s_cselect_b32 s58, 0x3f803f80, 0
+; CHECK-NEXT: s_and_b64 s[54:55], s[56:57], exec
+; CHECK-NEXT: v_accvgpr_write_b32 a0, s60
+; CHECK-NEXT: s_waitcnt lgkmcnt(0)
+; CHECK-NEXT: s_cselect_b32 s54, s11, 0
+; CHECK-NEXT: s_cselect_b32 s55, s10, 0
+; CHECK-NEXT: s_cselect_b32 s56, s9, 0
+; CHECK-NEXT: s_cselect_b32 s57, s8, 0
+; CHECK-NEXT: s_or_b32 s68, s1, s0
+; CHECK-NEXT: v_accvgpr_mov_b32 a1, a0
+; CHECK-NEXT: v_accvgpr_mov_b32 a2, a0
+; CHECK-NEXT: v_accvgpr_mov_b32 a3, a0
+; CHECK-NEXT: v_accvgpr_write_b32 a4, s58
+; CHECK-NEXT: s_and_b64 s[44:45], s[44:45], exec
+; CHECK-NEXT: s_waitcnt vmcnt(2)
+; CHECK-NEXT: v_cndmask_b32_e32 v19, 0, v19, vcc
+; CHECK-NEXT: v_cndmask_b32_e32 v18, 0, v18, vcc
+; CHECK-NEXT: v_cndmask_b32_e32 v17, 0, v17, vcc
+; CHECK-NEXT: v_cndmask_b32_e32 v16, 0, v16, vcc
+; CHECK-NEXT: v_accvgpr_mov_b32 a5, a4
+; CHECK-NEXT: v_accvgpr_mov_b32 a6, a4
+; CHECK-NEXT: v_accvgpr_mov_b32 a7, a4
+; CHECK-NEXT: s_cselect_b32 s69, s15, 0
+; CHECK-NEXT: s_cselect_b32 s70, s14, 0
+; CHECK-NEXT: s_cselect_b32 s71, s13, 0
+; CHECK-NEXT: s_cselect_b32 s16, s12, 0
+; CHECK-NEXT: s_and_b64 s[66:67], s[6:7], exec
+; CHECK-NEXT: s_cselect_b32 s17, s64, 0
+; CHECK-NEXT: s_cselect_b32 s18, s63, 0
+; CHECK-NEXT: s_cselect_b32 s19, s62, 0
+; CHECK-NEXT: s_load_dwordx4 s[28:31], s[34:35], 0x20
+; CHECK-NEXT: s_mov_b64 s[60:61], s[52:53]
+; CHECK-NEXT: s_mov_b64 s[52:53], 0
+; CHECK-NEXT: s_mov_b64 s[58:59], s[40:41]
+; CHECK-NEXT: s_mov_b32 s1, 1
+; CHECK-NEXT: s_mov_b64 s[44:45], -1
+; CHECK-NEXT: s_mov_b32 s62, 1
+; CHECK-NEXT: s_mov_b32 s63, 1
+; CHECK-NEXT: s_mov_b32 s64, 1
+; CHECK-NEXT: s_waitcnt vmcnt(1)
+; CHECK-NEXT: v_mfma_f32_32x32x16_bf16 v[0:15], a[8:11], a[8:11], v[0:15]
+; CHECK-NEXT: s_nop 11
+; CHECK-NEXT: scratch_store_dwordx4 off, v[0:3], off ; 16-byte Folded Spill
+; CHECK-NEXT: s_nop 0
+; CHECK-NEXT: scratch_store_dwordx4 off, v[4:7], off offset:16 ; 16-byte Folded Spill
+; CHECK-NEXT: scratch_store_dwordx4 off, v[8:11], off offset:32 ; 16-byte Folded Spill
+; CHECK-NEXT: scratch_store_dwordx4 off, v[12:15], off offset:48 ; 16-byte Folded Spill
+; CHECK-NEXT: scratch_load_dword v0, off, off offset:80 ; 4-byte Folded Reload
+; CHECK-NEXT: v_mfma_f32_32x32x16_bf16 a[16:31], a[8:11], a[8:11], a[16:31]
+; CHECK-NEXT: s_waitcnt vmcnt(0)
+; CHECK-NEXT: ds_write_b128 v0, a[0:3]
+; CHECK-NEXT: v_accvgpr_write_b32 a0, s57
+; CHECK-NEXT: v_accvgpr_write_b32 a1, s56
+; CHECK-NEXT: v_accvgpr_write_b32 a2, s55
+; CHECK-NEXT: v_accvgpr_write_b32 a3, s54
+; CHECK-NEXT: ds_write_b128 v30, a[4:7]
+; CHECK-NEXT: ds_write_b128 v20, a[0:3]
+; CHECK-NEXT: ds_write_b128 v30, v[16:19]
+; CHECK-NEXT: v_mov_b32_e32 v16, s16
+; CHECK-NEXT: v_mov_b32_e32 v17, s71
+; CHECK-NEXT: v_mov_b32_e32 v18, s70
+; CHECK-NEXT: v_mov_b32_e32 v19, s69
+; CHECK-NEXT: s_cselect_b32 s16, s65, 0
+; CHECK-NEXT: ds_write_b128 v30, v[16:19]
+; CHECK-NEXT: v_mov_b32_e32 v16, s19
+; CHECK-NEXT: v_mov_b32_e32 v17, s18
+; CHECK-NEXT: v_mov_b32_e32 v18, s17
+; CHECK-NEXT: v_mov_b32_e32 v19, s16
+; CHECK-NEXT: ds_write_b128 v20, v[16:19]
+; CHECK-NEXT: scratch_load_dwordx4 v[14:17], off, off ; 16-byte Folded Reload
+; CHECK-NEXT: scratch_load_dwordx4 v[18:21], off, off offset:16 ; 16-byte Folded Reload
+; CHECK-NEXT: scratch_load_dwordx4 v[22:25], off, off offset:32 ; 16-byte Folded Reload
+; CHECK-NEXT: scratch_load_dwordx4 v[26:29], off, off offset:48 ; 16-byte Folded Reload
+; CHECK-NEXT: v_readlane_b32 s18, v31, 4
+; CHECK-NEXT: s_and_b64 s[46:47], s[46:47], exec
+; CHECK-NEXT: v_readlane_b32 s19, v31, 5
+; CHECK-NEXT: s_cselect_b32 s16, 0x3f803f80, 0
+; CHECK-NEXT: s_and_b64 s[66:67], s[18:19], exec
+; CHECK-NEXT: v_mov_b32_e32 v0, s16
+; CHECK-NEXT: s_cselect_b32 s16, s33, 0
+; CHECK-NEXT: s_and_b64 s[48:49], s[48:49], exec
+; CHECK-NEXT: s_cselect_b32 s17, 0x3f803f80, 0
+; CHECK-NEXT: s_lshr_b32 s16, s16, 16
+; CHECK-NEXT: s_lshr_b32 s17, s17, 16
+; CHECK-NEXT: s_and_b64 s[50:51], s[50:51], exec
+; CHECK-NEXT: s_pack_ll_b32_b16 s16, s16, s17
+; CHECK-NEXT: s_waitcnt lgkmcnt(0)
+; CHECK-NEXT: s_cselect_b32 s17, s31, 0
+; CHECK-NEXT: v_mov_b32_e32 v1, v0
+; CHECK-NEXT: v_mov_b32_e32 v2, v0
+; CHECK-NEXT: v_mov_b32_e32 v3, v0
+; CHECK-NEXT: ds_write_b128 v30, v[0:3]
+; CHECK-NEXT: v_readlane_b32 s48, v31, 6
+; CHECK-NEXT: s_mov_b64 s[54:55], s[42:43]
+; CHECK-NEXT: s_mov_b64 s[56:57], s[38:39]
+; CHECK-NEXT: s_mov_b64 s[46:47], -1
+; CHECK-NEXT: s_mov_b32 s65, 1
+; CHECK-NEXT: v_readlane_b32 s49, v31, 7
+; CHECK-NEXT: s_mov_b32 s33, s23
+; CHECK-NEXT: s_mov_b64 s[50:51], s[36:37]
+; CHECK-NEXT: s_waitcnt vmcnt(0)
+; CHECK-NEXT: v_mfma_f32_32x32x16_bf16 v[14:29], a[8:11], a[8:11], v[14:29]
+; CHECK-NEXT: s_nop 11
+; CHECK-NEXT: v_accvgpr_write_b32 a0, v14
+; CHECK-NEXT: v_accvgpr_write_b32 a1, v15
+; CHECK-NEXT: v_accvgpr_write_b32 a2, v16
+; CHECK-NEXT: v_accvgpr_write_b32 a3, v17
+; CHECK-NEXT: v_accvgpr_write_b32 a4, v18
+; CHECK-NEXT: v_accvgpr_write_b32 a5, v19
+; CHECK-NEXT: v_accvgpr_write_b32 a6, v20
+; CHECK-NEXT: v_accvgpr_write_b32 a7, v21
+; CHECK-NEXT: v_accvgpr_write_b32 a8, v22
+; CHECK-NEXT: v_accvgpr_write_b32 a9, v23
+; CHECK-NEXT: v_accvgpr_write_b32 a10, v24
+; CHECK-NEXT: v_accvgpr_write_b32 a11, v25
+; CHECK-NEXT: v_accvgpr_write_b32 a12, v26
+; CHECK-NEXT: v_accvgpr_write_b32 a13, v27
+; CHECK-NEXT: v_accvgpr_write_b32 a14, v28
+; CHECK-NEXT: v_accvgpr_write_b32 a15, v29
+; CHECK-NEXT: v_mov_b32_e32 v16, s16
+; CHECK-NEXT: s_lshr_b32 s16, s17, 16
+; CHECK-NEXT: v_mov_b32_e32 v17, s16
+; CHECK-NEXT: ds_write_b64 v30, v[16:17]
+; CHECK-NEXT: v_mov_b32_e32 v16, s68
+; CHECK-NEXT: buffer_load_dwordx4 v[16:19], v16, s[24:27], 0 offen
+; CHECK-NEXT: s_and_b64 s[66:67], s[4:5], exec
+; CHECK-NEXT: s_cselect_b32 s16, 1, 0
+; CHECK-NEXT: s_cmp_lg_u32 s16, 1
+; CHECK-NEXT: s_cbranch_scc1 .LBB0_1
+; CHECK-NEXT: ; %bb.2: ; %bbl290
+; CHECK-NEXT: s_load_dwordx4 s[4:7], s[34:35], 0x80
+; CHECK-NEXT: s_and_b64 s[0:1], s[60:61], exec
+; CHECK-NEXT: s_cselect_b32 s0, 0x3f803f80, 0
+; CHECK-NEXT: v_mov_b32_e32 v4, s0
+; CHECK-NEXT: v_mov_b32_e32 v5, s0
+; CHECK-NEXT: v_mov_b32_e32 v6, s0
+; CHECK-NEXT: v_mov_b32_e32 v7, s0
+; CHECK-NEXT: v_mov_b32_e32 v0, 0
+; CHECK-NEXT: ds_write_b128 v0, v[4:7]
+; CHECK-NEXT: v_mov_b32_e32 v4, s3
+; CHECK-NEXT: s_waitcnt lgkmcnt(0)
+; CHECK-NEXT: v_mov_b64_e32 v[8:9], s[6:7]
+; CHECK-NEXT: v_mov_b64_e32 v[6:7], s[4:5]
+; CHECK-NEXT: ds_write_b128 v4, v[6:9]
+; CHECK-NEXT: v_mov_b64_e32 v[4:5], s[20:21]
+; CHECK-NEXT: v_mov_b64_e32 v[6:7], s[22:23]
+; CHECK-NEXT: ds_write_b128 v0, v[4:7]
+; CHECK-NEXT: v_mov_b32_e32 v4, s2
+; CHECK-NEXT: v_readlane_b32 s0, v31, 0
+; CHECK-NEXT: v_mov_b32_e32 v1, v0
+; CHECK-NEXT: v_mov_b32_e32 v2, v0
+; CHECK-NEXT: v_mov_b32_e32 v3, v0
+; CHECK-NEXT: v_readlane_b32 s1, v31, 1
+; CHECK-NEXT: v_readlane_b32 s2, v31, 2
+; CHECK-NEXT: v_readlane_b32 s3, v31, 3
+; CHECK-NEXT: ds_write_b128 v4, v[0:3]
+; CHECK-NEXT: v_mov_b64_e32 v[2:3], s[0:1]
+; CHECK-NEXT: v_mov_b64_e32 v[4:5], s[2:3]
+; CHECK-NEXT: ds_write_b128 v0, v[2:5]
+; CHECK-NEXT: s_endpgm
bbl:
br label %bbl97
bbl97: ; preds = %bbl97, %bbl
- %phi = phi i1 [ false, %bbl ], [ %and171, %bbl97 ]
- %phi98 = phi i1 [ false, %bbl ], [ %and175, %bbl97 ]
- %phi99 = phi i1 [ false, %bbl ], [ %and179, %bbl97 ]
- %phi101 = phi i1 [ false, %bbl ], [ %arg15, %bbl97 ]
- %phi102 = phi i1 [ false, %bbl ], [ %arg26, %bbl97 ]
- %phi103 = phi i1 [ false, %bbl ], [ %arg37, %bbl97 ]
- %phi104 = phi i1 [ false, %bbl ], [ %arg16, %bbl97 ]
- %phi106 = phi i32 [ 0, %bbl ], [ 1, %bbl97 ]
- %phi1071 = phi i32 [ 0, %bbl ], [ 1, %bbl97 ]
+ %phi = phi i1 [ false, %bbl ], [ %icmp, %bbl97 ]
+ %phi98 = phi i1 [ false, %bbl ], [ %phi103, %bbl97 ]
+ %phi99 = phi i1 [ false, %bbl ], [ %and256, %bbl97 ]
+ %phi101 = phi i1 [ false, %bbl ], [ true, %bbl97 ]
+ %phi104 = phi i1 [ false, %bbl ], [ true, %bbl97 ]
+ %phi105 = phi i32 [ 0, %bbl ], [ 1, %bbl97 ]
+ %icmp170 = phi i1 [ true, %bbl ], [ false, %bbl97 ]
+ %phi109 = phi <4 x i32> [ zeroinitializer, %bbl ], [ splat (i32 1), %bbl97 ]
%phi112 = phi <4 x i32> [ zeroinitializer, %bbl ], [ %call184, %bbl97 ]
- %phi113 = phi <4 x i32> [ zeroinitializer, %bbl ], [ splat (i32 1), %bbl97 ]
- %phi116 = phi i1 [ false, %bbl ], [ %and250, %bbl97 ]
- %phi117 = phi i1 [ false, %bbl ], [ %and256, %bbl97 ]
+ %phi117 = phi i1 [ false, %bbl ], [ %arg37, %bbl97 ]
%phi118 = phi i1 [ false, %bbl ], [ %arg38, %bbl97 ]
- %phi121 = phi i32 [ 0, %bbl ], [ 1, %bbl97 ]
- %phi122 = phi <8 x bfloat> [ zeroinitializer, %bbl ], [ %sitofp272, %bbl97 ]
+ %phi125 = phi <8 x bfloat> [ zeroinitializer, %bbl ], [ %sitofp, %bbl97 ]
+ %phi127 = phi <16 x float> [ zeroinitializer, %bbl ], [ %call282, %bbl97 ]
%phi128 = phi <16 x float> [ zeroinitializer, %bbl ], [ %call281, %bbl97 ]
- %bitcast162 = bitcast <4 x i32> %phi112 to <8 x bfloat>
- %select163 = select i1 %arg21, <8 x bfloat> %bitcast162, <8 x bfloat> zeroinitializer
+ %call154 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x16.bf16(<8 x bfloat> zeroinitializer, <8 x bfloat> zeroinitializer, <16 x float> %phi127, i32 0, i32 0, i32 0)
%select = select i1 %phi, <8 x bfloat> splat (bfloat 1.000000e+00), <8 x bfloat> zeroinitializer
- %getelementptr164 = getelementptr [2 x i8], ptr addrspace(3) null, i32 %phi1071
+ %getelementptr164 = getelementptr [2 x i8], ptr addrspace(3) null, i64 %0
store <8 x bfloat> %select, ptr addrspace(3) %getelementptr164, align 16
%select159 = select i1 %phi98, <8 x bfloat> splat (bfloat 1.000000e+00), <8 x bfloat> zeroinitializer
store <8 x bfloat> %select159, ptr addrspace(3) null, align 16
- %bitcast160 = bitcast <4 x i32> %phi113 to <8 x bfloat>
%select161 = select i1 %phi99, <8 x bfloat> %bitcast160, <8 x bfloat> zeroinitializer
store <8 x bfloat> %select161, ptr addrspace(3) %arg23, align 16
+ %bitcast162 = bitcast <4 x i32> %phi112 to <8 x bfloat>
+ %select163 = select i1 %arg21, <8 x bfloat> %bitcast162, <8 x bfloat> zeroinitializer
store <8 x bfloat> %select163, ptr addrspace(3) null, align 16
- %call1 = tail call i32 @llvm.amdgcn.workitem.id.x()
- %icmp = icmp slt i32 %call1, 1
- %icmp170 = icmp slt i32 %phi106, 1
- %and171 = and i1 %icmp, %icmp170
- %icmp73 = icmp slt i32 %call1, %arg6
- %and175 = and i1 %icmp73, %icmp170
- %and = and i32 %call1, 1020
- %icmp75 = icmp slt i32 %and, 1
- %and179 = and i1 %icmp75, %icmp170
- %call184 = tail call <4 x i32> @llvm.amdgcn.raw.ptr.buffer.load.v4i32(ptr addrspace(8) null, i32 0, i32 0, i32 0)
- %select219 = select i1 %phi104, <8 x bfloat> %bitcast218, <8 x bfloat> zeroinitializer
- %select213 = select i1 %phi101, <8 x bfloat> %sitofp, <8 x bfloat> zeroinitializer
- store <8 x bfloat> %select213, ptr addrspace(3) %arg23, align 16
- %select215 = select i1 %phi102, <8 x bfloat> %bitcast214, <8 x bfloat> zeroinitializer
- store <8 x bfloat> %select215, ptr addrspace(3) null, align 16
- %select217 = select i1 %phi103, <8 x bfloat> %bitcast216, <8 x bfloat> zeroinitializer
- store <8 x bfloat> %select217, ptr addrspace(3) %arg31, align 16
- store <8 x bfloat> %select219, ptr addrspace(3) %arg35, align 16
- %select227 = select i1 %phi118, <8 x bfloat> %phi123, <8 x bfloat> zeroinitializer
- %shufflevector230 = shufflevector <8 x bfloat> %select227, <8 x bfloat> zeroinitializer, <4 x i32> <i32 poison, i32 poison, i32 poison, i32 7>
- %select225 = select i1 %phi116, <8 x bfloat> %bitcast212, <8 x bfloat> zeroinitializer
+ %add174 = or i32 %phi105, %arg6
+ %call184 = tail call <4 x i32> @llvm.amdgcn.raw.ptr.buffer.load.v4i32(ptr addrspace(8) null, i32 %add174, i32 0, i32 0)
+ %select213 = select i1 %phi101, <8 x bfloat> %bitcast218, <8 x bfloat> zeroinitializer
+ store <8 x bfloat> %select213, ptr addrspace(3) null, align 16
+ %bitcast216 = bitcast <4 x i32> %phi109 to <8 x bfloat>
+ %select217 = select i1 %arg15, <8 x bfloat> %bitcast216, <8 x bfloat> zeroinitializer
+ store <8 x bfloat> %select217, ptr addrspace(3) %arg23, align 16
+ %select219 = select i1 %phi104, <8 x bfloat> splat (bfloat 1.000000e+00), <8 x bfloat> zeroinitializer
+ store <8 x bfloat> %select219, ptr addrspace(3) null, align 16
+ %select225 = select i1 %arg16, <8 x bfloat> %phi125, <8 x bfloat> zeroinitializer
%select226 = select i1 %phi117, <8 x bfloat> splat (bfloat 1.000000e+00), <8 x bfloat> zeroinitializer
%shufflevector229 = shufflevector <8 x bfloat> %select225, <8 x bfloat> %select226, <4 x i32> <i32 7, i32 15, i32 poison, i32 poison>
+ %1 = shufflevector <8 x bfloat> %phi123, <8 x bfloat> zeroinitializer, <4 x i32> <i32 poison, i32 poison, i32 poison, i32 7>
+ %shufflevector230 = select i1 %phi118, <4 x bfloat> %1, <4 x bfloat> zeroinitializer
%shufflevector231 = shufflevector <4 x bfloat> %shufflevector229, <4 x bfloat> %shufflevector230, <4 x i32> <i32 0, i32 1, i32 7, i32 poison>
- %select228 = select i1 %arg27, <8 x bfloat> %phi122, <8 x bfloat> zeroinitializer
- %shufflevector232 = shufflevector <8 x bfloat> %select228, <8 x bfloat> zeroinitializer, <4 x i32> <i32 poison, i32 poison, i32 poison, i32 7>
- %shufflevector233 = shufflevector <4 x bfloat> %shufflevector231, <4 x bfloat> %shufflevector232, <4 x i32> <i32 0, i32 1, i32 2, i32 7>
- store <4 x bfloat> %shufflevector233, ptr addrspace(3) null, align 8
- %icmp249 = icmp slt i32 %phi121, %arg6
- %and250 = and i1 %phi119, %icmp249
- %and64 = and i32 %call1, 1
- %icmp783 = icmp slt i32 %and64, 1
- %icmp255 = icmp slt i32 %phi121, %arg36
- %and256 = and i1 %icmp783, %icmp255
- %sitofp272 = sitofp <8 x i8> %bitcast252 to <8 x bfloat>
- %call153 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x16.bf16(<8 x bfloat> splat (bfloat 1.000000e+00), <8 x bfloat> zeroinitializer, <16 x float> %phi128, i32 0, i32 0, i32 0)
- %call281 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x16.bf16(<8 x bfloat> zeroinitializer, <8 x bfloat> zeroinitializer, <16 x float> %call153, i32 0, i32 0, i32 0)
- tail call void @llvm.amdgcn.sched.barrier(i32 0)
+ store <4 x bfloat> %shufflevector231, ptr addrspace(3) null, align 8
+ %call281 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x16.bf16(<8 x bfloat> zeroinitializer, <8 x bfloat> zeroinitializer, <16 x float> %phi128, i32 0, i32 0, i32 0)
+ %call282 = tail call <16 x float> @llvm.amdgcn.mfma.f32.32x32x16.bf16(<8 x bfloat> zeroinitializer, <8 x bfloat> zeroinitializer, <16 x float> %call154, i32 0, i32 0, i32 0)
br i1 %arg15, label %bbl97, label %bbl290
bbl290: ; preds = %bbl97
- %select292 = select i1 %and171, <8 x bfloat> splat (bfloat 1.000000e+00), <8 x bfloat> zeroinitializer
+ %select292 = select i1 %icmp170, <8 x bfloat> splat (bfloat 1.000000e+00), <8 x bfloat> zeroinitializer
store <8 x bfloat> %select292, ptr addrspace(3) null, align 16
- %select294 = select i1 %and175, <8 x bfloat> splat (bfloat 1.000000e+00), <8 x bfloat> zeroinitializer
- store <8 x bfloat> %select294, ptr addrspace(3) %arg46, align 16
- %select296 = select i1 %and179, <8 x bfloat> splat (bfloat 1.000000e+00), <8 x bfloat> zeroinitializer
- store <8 x bfloat> %select296, ptr addrspace(3) null, align 16
- %getelementptr301 = getelementptr [2 x i8], ptr addrspace(3) null, i32 %arg11
- store <8 x bfloat> zeroinitializer, ptr addrspace(3) %getelementptr301, align 16
- %select314 = select i1 %and256, <8 x bfloat> splat (bfloat 1.000000e+00), <8 x bfloat> zeroinitializer
- %shufflevector316 = shufflevector <8 x bfloat> zeroinitializer, <8 x bfloat> %select314, <4 x i32> <i32 0, i32 8, i32 poison, i32 poison>
- %shufflevector318 = shufflevector <4 x bfloat> %shufflevector316, <4 x bfloat> zeroinitializer, <4 x i32> <i32 0, i32 1, i32 4, i32 poison>
- %shufflevector319 = shufflevector <4 x bfloat> %shufflevector318, <4 x bfloat> zeroinitializer, <4 x i32> <i32 0, i32 1, i32 2, i32 4>
+ store <8 x bfloat> %select294, ptr addrspace(3) %arg33, align 16
+ store <8 x bfloat> %sitofp, ptr addrspace(3) null, align 16
+ store <8 x bfloat> zeroinitializer, ptr addrspace(3) %arg23, align 16
+ store <8 x bfloat> %phi123, ptr addrspace(3) null, align 16
ret void
}
-attributes #0 = { nocallback nofree nosync nounwind speculatable willreturn memory(none) }
-attributes #1 = { nocallback nofree nosync nounwind willreturn memory(argmem: read) }
-attributes #2 = { convergent nocallback nofree nounwind willreturn }
-attributes #3 = { convergent nocallback nocreateundeforpoison nofree nosync nounwind willreturn memory(none) }
-attributes #4 = { "amdgpu-num-vgpr"="24" }
+attributes #0 = { "amdgpu-waves-per-eu"="8,8" }
More information about the llvm-commits
mailing list