[llvm] accc841 - [AMDGPU] Refresh RegisterClassInfo after reserving WWM registers (#221849)
via llvm-commits
llvm-commits at lists.llvm.org
Tue Sep 8 13:48:05 PDT 2026
Author: Michael Halkenhäuser
Date: 2026-09-08T22:48:00+02:00
New Revision: accc84148aae88d54c11403d04929afaccecf4d6
URL: https://github.com/llvm/llvm-project/commit/accc84148aae88d54c11403d04929afaccecf4d6
DIFF: https://github.com/llvm/llvm-project/commit/accc84148aae88d54c11403d04929afaccecf4d6.diff
LOG: [AMDGPU] Refresh RegisterClassInfo after reserving WWM registers (#221849)
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.
Added:
llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-reserved-wwm-vgpr.ll
Modified:
llvm/lib/Target/AMDGPU/AMDGPUReserveWWMRegs.cpp
Removed:
################################################################################
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;
}
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..2f9c62fd283d0
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/rewrite-vgpr-mfma-to-agpr-reserved-wwm-vgpr.ll
@@ -0,0 +1,379 @@
+; 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
+
+; "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.
+
+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, 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 ], [ %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 ]
+ %phi117 = phi i1 [ false, %bbl ], [ %arg37, %bbl97 ]
+ %phi118 = phi i1 [ false, %bbl ], [ %arg38, %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 ]
+ %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, 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
+ %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
+ %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>
+ 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 %icmp170, <8 x bfloat> splat (bfloat 1.000000e+00), <8 x bfloat> zeroinitializer
+ store <8 x bfloat> %select292, ptr addrspace(3) null, align 16
+ 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 = { "amdgpu-waves-per-eu"="8,8" }
More information about the llvm-commits
mailing list