From 25ad7cd743f2cd93f535a01e1df6b2e449fa92a6 Mon Sep 17 00:00:00 2001 From: Matt Arsenault Date: Sun, 5 Oct 2025 11:21:02 +0900 Subject: [PATCH 1/2] AMDGPU: Fix mfma agpr allocation failures with -O0 Previously we were getting lucky on cases that can use AV registers with the normal optimization pipeline. I do not understand what the check against getAddressableNumArchVGPRs was doing here. This logic needs to be consistent with getMaxNumVectorRegs, as that is what getReservedRegs to determine the AGPR budget. In the future we should directly check the minimum AGPR budget, and individual selection patterns need to know the minimum budget required for them. --- .../Target/AMDGPU/SIMachineFunctionInfo.cpp | 4 +- .../AMDGPU/smfmac_alloc_failure_no_agpr_O0.ll | 119 ++++++++++++++++++ 2 files changed, 120 insertions(+), 3 deletions(-) create mode 100644 llvm/test/CodeGen/AMDGPU/smfmac_alloc_failure_no_agpr_O0.ll diff --git a/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.cpp b/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.cpp index 908d856d386f5..ec4e4f9ea9a3a 100644 --- a/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.cpp +++ b/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.cpp @@ -85,9 +85,7 @@ SIMachineFunctionInfo::SIMachineFunctionInfo(const Function &F, if (ST.hasGFX90AInsts()) { // FIXME: MayNeedAGPRs is a misnomer for how this is used. MFMA selection // should be separated from availability of AGPRs - if (MFMAVGPRForm || - (ST.getMaxNumVGPRs(F) <= ST.getAddressableNumArchVGPRs() && - !mayUseAGPRs(F))) + if (!mayUseAGPRs(F)) MayNeedAGPRs = false; // We will select all MAI with VGPR operands. } diff --git a/llvm/test/CodeGen/AMDGPU/smfmac_alloc_failure_no_agpr_O0.ll b/llvm/test/CodeGen/AMDGPU/smfmac_alloc_failure_no_agpr_O0.ll new file mode 100644 index 0000000000000..ba0fdc689b4ff --- /dev/null +++ b/llvm/test/CodeGen/AMDGPU/smfmac_alloc_failure_no_agpr_O0.ll @@ -0,0 +1,119 @@ +; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6 +; RUN: llc -O0 -mtriple=amdgcn -mcpu=gfx950 -amdgpu-mfma-vgpr-form=0 < %s | FileCheck %s +; RUN: llc -O0 -mtriple=amdgcn -mcpu=gfx950 -amdgpu-mfma-vgpr-form=1 < %s | FileCheck %s + +declare <16 x float> @llvm.amdgcn.smfmac.f32.32x32x32.f16(<8 x half>, <16 x half>, <16 x float>, i32, i32 immarg, i32 immarg) + +define amdgpu_kernel void @test_smfmac_f32_32x32x32_f16__vgpr(ptr addrspace(1) %arg, <8 x half> %a, <16 x half> %b, i32 %idx) #0 { +; CHECK-LABEL: test_smfmac_f32_32x32x32_f16__vgpr: +; CHECK: ; %bb.0: ; %bb +; CHECK-NEXT: s_mov_b64 s[2:3], s[4:5] +; CHECK-NEXT: v_mov_b32_e32 v1, v0 +; CHECK-NEXT: v_mov_b32_e32 v0, 0 +; CHECK-NEXT: s_load_dwordx2 s[0:1], s[2:3], 0x24 +; CHECK-NEXT: s_load_dwordx4 s[12:15], s[2:3], 0x34 +; CHECK-NEXT: s_load_dwordx8 s[4:11], s[2:3], 0x44 +; CHECK-NEXT: s_nop 0 +; CHECK-NEXT: s_load_dword s2, s[2:3], 0x64 +; CHECK-NEXT: s_mov_b32 s3, 0x3ff +; CHECK-NEXT: v_and_b32_e64 v1, v1, s3 +; CHECK-NEXT: s_mov_b32 s3, 6 +; CHECK-NEXT: v_lshlrev_b32_e64 v8, s3, v1 +; CHECK-NEXT: s_waitcnt lgkmcnt(0) +; CHECK-NEXT: global_load_dwordx4 v[4:7], v8, s[0:1] offset:48 +; CHECK-NEXT: s_waitcnt vmcnt(0) +; CHECK-NEXT: v_mov_b32_e32 v1, v7 +; CHECK-NEXT: v_mov_b32_e32 v2, v6 +; CHECK-NEXT: v_mov_b32_e32 v3, v5 +; CHECK-NEXT: ; kill: def $vgpr4 killed $vgpr4 killed $vgpr4_vgpr5_vgpr6_vgpr7 killed $exec +; CHECK-NEXT: global_load_dwordx4 v[10:13], v8, s[0:1] offset:32 +; CHECK-NEXT: s_waitcnt vmcnt(0) +; CHECK-NEXT: v_mov_b32_e32 v5, v13 +; CHECK-NEXT: v_mov_b32_e32 v6, v12 +; CHECK-NEXT: v_mov_b32_e32 v7, v11 +; CHECK-NEXT: v_mov_b32_e32 v24, v10 +; CHECK-NEXT: global_load_dwordx4 v[10:13], v8, s[0:1] offset:16 +; CHECK-NEXT: s_waitcnt vmcnt(0) +; CHECK-NEXT: v_mov_b32_e32 v25, v13 +; CHECK-NEXT: v_mov_b32_e32 v26, v12 +; CHECK-NEXT: v_mov_b32_e32 v27, v11 +; CHECK-NEXT: v_mov_b32_e32 v28, v10 +; CHECK-NEXT: global_load_dwordx4 v[8:11], v8, s[0:1] +; CHECK-NEXT: s_waitcnt vmcnt(0) +; CHECK-NEXT: v_mov_b32_e32 v29, v11 +; CHECK-NEXT: v_mov_b32_e32 v30, v10 +; CHECK-NEXT: v_mov_b32_e32 v31, v9 +; CHECK-NEXT: ; kill: def $vgpr8 killed $vgpr8 killed $vgpr8_vgpr9_vgpr10_vgpr11 killed $exec +; CHECK-NEXT: ; kill: def $vgpr8 killed $vgpr8 def $vgpr8_vgpr9_vgpr10_vgpr11_vgpr12_vgpr13_vgpr14_vgpr15_vgpr16_vgpr17_vgpr18_vgpr19_vgpr20_vgpr21_vgpr22_vgpr23 killed $exec +; CHECK-NEXT: v_mov_b32_e32 v9, v31 +; CHECK-NEXT: v_mov_b32_e32 v10, v30 +; CHECK-NEXT: v_mov_b32_e32 v11, v29 +; CHECK-NEXT: v_mov_b32_e32 v12, v28 +; CHECK-NEXT: v_mov_b32_e32 v13, v27 +; CHECK-NEXT: v_mov_b32_e32 v14, v26 +; CHECK-NEXT: v_mov_b32_e32 v15, v25 +; CHECK-NEXT: v_mov_b32_e32 v16, v24 +; CHECK-NEXT: v_mov_b32_e32 v17, v7 +; CHECK-NEXT: v_mov_b32_e32 v18, v6 +; CHECK-NEXT: v_mov_b32_e32 v19, v5 +; CHECK-NEXT: v_mov_b32_e32 v20, v4 +; CHECK-NEXT: v_mov_b32_e32 v21, v3 +; CHECK-NEXT: v_mov_b32_e32 v22, v2 +; CHECK-NEXT: v_mov_b32_e32 v23, v1 +; CHECK-NEXT: v_mov_b64_e32 v[2:3], s[12:13] +; CHECK-NEXT: v_mov_b64_e32 v[4:5], s[14:15] +; CHECK-NEXT: v_mov_b64_e32 v[30:31], s[10:11] +; CHECK-NEXT: v_mov_b64_e32 v[28:29], s[8:9] +; CHECK-NEXT: v_mov_b64_e32 v[26:27], s[6:7] +; CHECK-NEXT: v_mov_b64_e32 v[24:25], s[4:5] +; CHECK-NEXT: v_mov_b32_e32 v1, s2 +; CHECK-NEXT: s_nop 1 +; CHECK-NEXT: v_smfmac_f32_32x32x32_f16 v[8:23], v[2:5], v[24:31], v1 cbsz:1 abid:2 +; CHECK-NEXT: s_nop 11 +; CHECK-NEXT: v_mov_b32_e32 v1, v23 +; CHECK-NEXT: v_mov_b32_e32 v6, v22 +; CHECK-NEXT: v_mov_b32_e32 v7, v21 +; CHECK-NEXT: v_mov_b32_e32 v2, v20 +; CHECK-NEXT: ; kill: def $vgpr2 killed $vgpr2 def $vgpr2_vgpr3_vgpr4_vgpr5 killed $exec +; CHECK-NEXT: v_mov_b32_e32 v3, v7 +; CHECK-NEXT: v_mov_b32_e32 v4, v6 +; CHECK-NEXT: v_mov_b32_e32 v5, v1 +; CHECK-NEXT: global_store_dwordx4 v0, v[2:5], s[0:1] offset:48 +; CHECK-NEXT: v_mov_b32_e32 v1, v19 +; CHECK-NEXT: v_mov_b32_e32 v6, v18 +; CHECK-NEXT: v_mov_b32_e32 v7, v17 +; CHECK-NEXT: v_mov_b32_e32 v2, v16 +; CHECK-NEXT: ; kill: def $vgpr2 killed $vgpr2 def $vgpr2_vgpr3_vgpr4_vgpr5 killed $exec +; CHECK-NEXT: v_mov_b32_e32 v3, v7 +; CHECK-NEXT: v_mov_b32_e32 v4, v6 +; CHECK-NEXT: v_mov_b32_e32 v5, v1 +; CHECK-NEXT: global_store_dwordx4 v0, v[2:5], s[0:1] offset:32 +; CHECK-NEXT: v_mov_b32_e32 v1, v15 +; CHECK-NEXT: v_mov_b32_e32 v6, v14 +; CHECK-NEXT: v_mov_b32_e32 v7, v13 +; CHECK-NEXT: v_mov_b32_e32 v2, v12 +; CHECK-NEXT: ; kill: def $vgpr2 killed $vgpr2 def $vgpr2_vgpr3_vgpr4_vgpr5 killed $exec +; CHECK-NEXT: v_mov_b32_e32 v3, v7 +; CHECK-NEXT: v_mov_b32_e32 v4, v6 +; CHECK-NEXT: v_mov_b32_e32 v5, v1 +; CHECK-NEXT: global_store_dwordx4 v0, v[2:5], s[0:1] offset:16 +; CHECK-NEXT: v_mov_b32_e32 v1, v11 +; CHECK-NEXT: v_mov_b32_e32 v6, v10 +; CHECK-NEXT: v_mov_b32_e32 v7, v9 +; CHECK-NEXT: v_mov_b32_e32 v2, v8 +; CHECK-NEXT: ; kill: def $vgpr2 killed $vgpr2 def $vgpr2_vgpr3_vgpr4_vgpr5 killed $exec +; CHECK-NEXT: v_mov_b32_e32 v3, v7 +; CHECK-NEXT: v_mov_b32_e32 v4, v6 +; CHECK-NEXT: v_mov_b32_e32 v5, v1 +; CHECK-NEXT: global_store_dwordx4 v0, v[2:5], s[0:1] +; CHECK-NEXT: s_endpgm +bb: + %id = call i32 @llvm.amdgcn.workitem.id.x() + %gep = getelementptr <16 x float>, ptr addrspace(1) %arg, i32 %id + %in.1 = load <16 x float>, ptr addrspace(1) %gep + %mai.1 = tail call <16 x float> @llvm.amdgcn.smfmac.f32.32x32x32.f16(<8 x half> %a, <16 x half> %b, <16 x float> %in.1, i32 %idx, i32 1, i32 2) + store <16 x float> %mai.1, ptr addrspace(1) %arg + ret void +} + +attributes #0 = { "amdgpu-flat-work-group-size"="1,256" "amdgpu-agpr-alloc"="0,0" } From 764d91d76d9cc0ca1c4d2daa54e80d61bd273cad Mon Sep 17 00:00:00 2001 From: Matt Arsenault Date: Sun, 5 Oct 2025 23:52:47 +0900 Subject: [PATCH 2/2] AMDGPU: Track minNumAGPRs in MFI instead of mayUseAGPRs Start accounting for the number of AGPRs required to perform the allocation. Refine the selection predicates to check this number is available, and default to selecting the VGPR case if there aren't enough. This avoids register allocation failures for the largest MFMAs with the default register budget. --- .../Target/AMDGPU/AMDGPURegisterBankInfo.cpp | 14 +- llvm/lib/Target/AMDGPU/SIISelLowering.cpp | 3 +- .../Target/AMDGPU/SIMachineFunctionInfo.cpp | 20 +- .../lib/Target/AMDGPU/SIMachineFunctionInfo.h | 10 +- llvm/lib/Target/AMDGPU/VOP3PInstructions.td | 70 +++--- .../CodeGen/AMDGPU/llvm.amdgcn.mfma.form.ll | 201 ++++++++++++++++++ 6 files changed, 264 insertions(+), 54 deletions(-) diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp index 848d9a5a9eb98..557d87f079beb 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp @@ -5043,6 +5043,9 @@ AMDGPURegisterBankInfo::getInstrMapping(const MachineInstr &MI) const { case Intrinsic::amdgcn_mfma_i32_16x16x64_i8: case Intrinsic::amdgcn_mfma_i32_32x32x32_i8: case Intrinsic::amdgcn_mfma_f32_16x16x32_bf16: { + unsigned DstSize = MRI.getType(MI.getOperand(0).getReg()).getSizeInBits(); + unsigned MinNumRegsRequired = DstSize / 32; + // Default for MAI intrinsics. // srcC can also be an immediate which can be folded later. // FIXME: Should we eventually add an alternative mapping with AGPR src @@ -5051,29 +5054,32 @@ AMDGPURegisterBankInfo::getInstrMapping(const MachineInstr &MI) const { // vdst, srcA, srcB, srcC const SIMachineFunctionInfo *Info = MF.getInfo(); OpdsMapping[0] = - Info->mayNeedAGPRs() + Info->getMinNumAGPRs() >= MinNumRegsRequired ? getAGPROpMapping(MI.getOperand(0).getReg(), MRI, *TRI) : getVGPROpMapping(MI.getOperand(0).getReg(), MRI, *TRI); OpdsMapping[2] = getVGPROpMapping(MI.getOperand(2).getReg(), MRI, *TRI); OpdsMapping[3] = getVGPROpMapping(MI.getOperand(3).getReg(), MRI, *TRI); OpdsMapping[4] = - Info->mayNeedAGPRs() + Info->getMinNumAGPRs() >= MinNumRegsRequired ? getAGPROpMapping(MI.getOperand(4).getReg(), MRI, *TRI) : getVGPROpMapping(MI.getOperand(4).getReg(), MRI, *TRI); break; } case Intrinsic::amdgcn_mfma_scale_f32_16x16x128_f8f6f4: case Intrinsic::amdgcn_mfma_scale_f32_32x32x64_f8f6f4: { + unsigned DstSize = MRI.getType(MI.getOperand(0).getReg()).getSizeInBits(); + unsigned MinNumRegsRequired = DstSize / 32; + const SIMachineFunctionInfo *Info = MF.getInfo(); OpdsMapping[0] = - Info->mayNeedAGPRs() + Info->getMinNumAGPRs() >= MinNumRegsRequired ? getAGPROpMapping(MI.getOperand(0).getReg(), MRI, *TRI) : getVGPROpMapping(MI.getOperand(0).getReg(), MRI, *TRI); OpdsMapping[2] = getVGPROpMapping(MI.getOperand(2).getReg(), MRI, *TRI); OpdsMapping[3] = getVGPROpMapping(MI.getOperand(3).getReg(), MRI, *TRI); OpdsMapping[4] = - Info->mayNeedAGPRs() + Info->getMinNumAGPRs() >= MinNumRegsRequired ? getAGPROpMapping(MI.getOperand(4).getReg(), MRI, *TRI) : getVGPROpMapping(MI.getOperand(4).getReg(), MRI, *TRI); diff --git a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp index e2334577884b7..24a54e7c6e960 100644 --- a/llvm/lib/Target/AMDGPU/SIISelLowering.cpp +++ b/llvm/lib/Target/AMDGPU/SIISelLowering.cpp @@ -17357,7 +17357,8 @@ void SITargetLowering::AdjustInstrPostInstrSelection(MachineInstr &MI, // use between vgpr and agpr as agpr tuples tend to be big. if (!MI.getDesc().operands().empty()) { unsigned Opc = MI.getOpcode(); - bool HasAGPRs = Info->mayNeedAGPRs(); + bool HasAGPRs = + !Subtarget->hasGFX90AInsts() || Info->getMinNumAGPRs() != 0; const SIRegisterInfo *TRI = Subtarget->getRegisterInfo(); int16_t Src2Idx = AMDGPU::getNamedOperandIdx(Opc, AMDGPU::OpName::src2); for (auto I : diff --git a/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.cpp b/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.cpp index ec4e4f9ea9a3a..b398db4f7caff 100644 --- a/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.cpp +++ b/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.cpp @@ -33,17 +33,20 @@ using namespace llvm; // optimal RC for Opc and Dest of MFMA. In particular, there are high RP cases // where it is better to produce the VGPR form (e.g. if there are VGPR users // of the MFMA result). -static cl::opt MFMAVGPRForm( - "amdgpu-mfma-vgpr-form", cl::Hidden, +static cl::opt MFMAVGPRFormOpt( + "amdgpu-mfma-vgpr-form", cl::desc("Whether to force use VGPR for Opc and Dest of MFMA. If " "unspecified, default to compiler heuristics"), - cl::init(false)); + cl::location(SIMachineFunctionInfo::MFMAVGPRForm), cl::init(false), + cl::Hidden); const GCNTargetMachine &getTM(const GCNSubtarget *STI) { const SITargetLowering *TLI = STI->getTargetLowering(); return static_cast(TLI->getTargetMachine()); } +bool SIMachineFunctionInfo::MFMAVGPRForm = false; + SIMachineFunctionInfo::SIMachineFunctionInfo(const Function &F, const GCNSubtarget *STI) : AMDGPUMachineFunction(F, *STI), Mode(F, *STI), GWSResourcePSV(getTM(STI)), @@ -81,12 +84,13 @@ SIMachineFunctionInfo::SIMachineFunctionInfo(const Function &F, PSInputAddr = AMDGPU::getInitialPSInputAddr(F); } - MayNeedAGPRs = ST.hasMAIInsts(); if (ST.hasGFX90AInsts()) { - // FIXME: MayNeedAGPRs is a misnomer for how this is used. MFMA selection - // should be separated from availability of AGPRs - if (!mayUseAGPRs(F)) - MayNeedAGPRs = false; // We will select all MAI with VGPR operands. + // FIXME: Extract logic out of getMaxNumVectorRegs; we need to apply the + // allocation granule and clamping. + auto [MinNumAGPRAttr, MaxNumAGPRAttr] = + AMDGPU::getIntegerPairAttribute(F, "amdgpu-agpr-alloc", {~0u, ~0u}, + /*OnlyFirstRequired=*/true); + MinNumAGPRs = MinNumAGPRAttr; } if (AMDGPU::isChainCC(CC)) { diff --git a/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h b/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h index 45606153db58e..b7dbb5994ee41 100644 --- a/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h +++ b/llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h @@ -509,7 +509,9 @@ class SIMachineFunctionInfo final : public AMDGPUMachineFunction, // user arguments. This is an offset from the KernargSegmentPtr. bool ImplicitArgPtr : 1; - bool MayNeedAGPRs : 1; + /// Minimum number of AGPRs required to allocate in the function. Only + /// relevant for gfx90a-gfx950. For gfx908, this should be infinite. + unsigned MinNumAGPRs = ~0u; // The hard-wired high half of the address of the global information table // for AMDPAL OS type. 0xffffffff represents no hard-wired high half, since @@ -537,6 +539,8 @@ class SIMachineFunctionInfo final : public AMDGPUMachineFunction, void MRI_NoteCloneVirtualRegister(Register NewReg, Register SrcReg) override; public: + static bool MFMAVGPRForm; + struct VGPRSpillToAGPR { SmallVector Lanes; bool FullyAllocated = false; @@ -1196,9 +1200,7 @@ class SIMachineFunctionInfo final : public AMDGPUMachineFunction, unsigned getMaxMemoryClusterDWords() const { return MaxMemoryClusterDWords; } - bool mayNeedAGPRs() const { - return MayNeedAGPRs; - } + unsigned getMinNumAGPRs() const { return MinNumAGPRs; } // \returns true if a function has a use of AGPRs via inline asm or // has a call which may use it. diff --git a/llvm/lib/Target/AMDGPU/VOP3PInstructions.td b/llvm/lib/Target/AMDGPU/VOP3PInstructions.td index 5daf860d540ca..3a0cc35dcd4ff 100644 --- a/llvm/lib/Target/AMDGPU/VOP3PInstructions.td +++ b/llvm/lib/Target/AMDGPU/VOP3PInstructions.td @@ -67,7 +67,7 @@ class VOP3P_Mix_Profile : VOP3P_Mix_Profile { let IsTrue16 = 1; - let IsRealTrue16 = 1; + let IsRealTrue16 = 1; let DstRC64 = getVALUDstForVT.ret; } @@ -950,7 +950,7 @@ class MFMA_F8F6F4_WithSizeTable_Helper : } // Currently assumes scaled instructions never have abid -class MAIFrag : PatFrag < +class MAIFrag : PatFrag < !if(Scaled, (ops node:$src0, node:$src1, node:$src2, node:$cbsz, node:$blgp, node:$src0_modifiers, node:$scale_src0, node:$src1_modifiers, node:$scale_src1), @@ -959,37 +959,30 @@ class MAIFrag; - -defvar MayNeedAGPRs = [{ - return MF->getInfo()->mayNeedAGPRs(); -}]; - -defvar MayNeedAGPRs_gisel = [{ - return MF.getInfo()->mayNeedAGPRs(); -}]; + (Op $src0, $src1, $src2, $cbsz, $blgp)))>; -defvar MayNotNeedAGPRs = [{ - return !MF->getInfo()->mayNeedAGPRs(); -}]; +class CanUseAGPR_MAI { + code PredicateCode = [{ + return !Subtarget->hasGFX90AInsts() || + (!SIMachineFunctionInfo::MFMAVGPRForm && + MF->getInfo()->getMinNumAGPRs() >= + }] # !srl(vt.Size, 5) # ");"; -defvar MayNotNeedAGPRs_gisel = [{ - return !MF.getInfo()->mayNeedAGPRs(); -}]; + code GISelPredicateCode = [{ + return !Subtarget->hasGFX90AInsts() || + (!SIMachineFunctionInfo::MFMAVGPRForm && + MF.getInfo()->getMinNumAGPRs() >= + }] # !srl(vt.Size, 5) # ");"; +} -class AgprMAIFrag : - MAIFrag { - let GISelPredicateCode = MayNeedAGPRs_gisel; -} + MAIFrag, + CanUseAGPR_MAI; class VgprMAIFrag : - MAIFrag { - let GISelPredicateCode = MayNotNeedAGPRs_gisel; -} + bit Scaled = false> : + MAIFrag; let isAsCheapAsAMove = 1, isReMaterializable = 1 in { defm V_ACCVGPR_READ_B32 : VOP3Inst<"v_accvgpr_read_b32", VOPProfileAccRead>; @@ -1037,16 +1030,19 @@ multiclass MAIInst { defvar NoDstOverlap = !cast("VOPProfileMAI_" # P).NoDstOverlap; + defvar ProfileAGPR = !cast("VOPProfileMAI_" # P); + defvar ProfileVGPR = !cast("VOPProfileMAI_" # P # "_VCD"); + let isConvergent = 1, mayRaiseFPException = 0, ReadsModeReg = 1 in { // FP32 denorm mode is respected, rounding mode is not. Exceptions are not supported. let Constraints = !if(NoDstOverlap, "@earlyclobber $vdst", "") in { - def _e64 : MAIInst("VOPProfileMAI_" # P), - !if(!or(NoDstOverlap, !eq(node, null_frag)), null_frag, AgprMAIFrag), Scaled>, + def _e64 : MAIInst), Scaled>, MFMATable<0, "AGPR", NAME # "_e64">; let OtherPredicates = [isGFX90APlus], Mnemonic = OpName in - def _vgprcd_e64 : MAIInst("VOPProfileMAI_" # P # "_VCD"), + def _vgprcd_e64 : MAIInst), Scaled>, MFMATable<0, "VGPR", NAME # "_vgprcd_e64", NAME # "_e64">; } @@ -1055,12 +1051,12 @@ multiclass MAIInst("VOPProfileMAI_" # P), - !if(!eq(node, null_frag), null_frag, AgprMAIFrag), Scaled>, + def "_mac_e64" : MAIInst), Scaled>, MFMATable<1, "AGPR", NAME # "_e64", NAME # "_mac_e64">; let OtherPredicates = [isGFX90APlus] in - def _mac_vgprcd_e64 : MAIInst("VOPProfileMAI_" # P # "_VCD"), + def _mac_vgprcd_e64 : MAIInst), Scaled>, MFMATable<1, "VGPR", NAME # "_vgprcd_e64", NAME # "_mac_e64">; } @@ -1074,11 +1070,11 @@ multiclass ScaledMAIInst_mc(!cast(UnscaledOpName#"_e64").Pfl).NoDstOverlap; + defvar Profile = !cast(!cast(UnscaledOpName#"_e64").Pfl); + defvar NoDstOverlap = Profile.NoDstOverlap; def _e64 : ScaledMAIInst(UnscaledOpName#"_e64"), !if(NoDstOverlap, null_frag, AgprMAIFrag)>, + !cast(UnscaledOpName#"_e64"), !if(NoDstOverlap, null_frag, AgprMAIFrag)>, MFMATable<0, "AGPR", NAME # "_e64">; def _vgprcd_e64 : ScaledMAIInst(UnscaledOpName # "_mac_e64"), AgprMAIFrag>, + !cast(UnscaledOpName # "_mac_e64"), AgprMAIFrag>, MFMATable<1, "AGPR", NAME # "_e64">; def _mac_vgprcd_e64 : ScaledMAIInst @request_no_agpr(<8 x half> %arg0, <8 x half> %arg1, <4 x flo ret <4 x float> %result } +; Make sure this selects the VGPR form, if AGPRs available, but not +; enough. +define amdgpu_kernel void @not_enough_agprs(ptr addrspace(1) %arg) #2 { +; HEURRC-LABEL: not_enough_agprs: +; HEURRC: ; %bb.0: ; %bb +; HEURRC-NEXT: s_load_dwordx2 s[34:35], s[4:5], 0x24 +; HEURRC-NEXT: v_mov_b32_e32 v33, 1.0 +; HEURRC-NEXT: v_mov_b32_e32 v34, 2.0 +; HEURRC-NEXT: v_mov_b32_e32 v32, 0 +; HEURRC-NEXT: s_waitcnt lgkmcnt(0) +; HEURRC-NEXT: s_load_dwordx16 s[16:31], s[34:35], 0x0 +; HEURRC-NEXT: s_load_dwordx16 s[0:15], s[34:35], 0x40 +; HEURRC-NEXT: s_waitcnt lgkmcnt(0) +; HEURRC-NEXT: v_mov_b32_e32 v0, s16 +; HEURRC-NEXT: v_mov_b32_e32 v1, s17 +; HEURRC-NEXT: v_mov_b32_e32 v2, s18 +; HEURRC-NEXT: v_mov_b32_e32 v3, s19 +; HEURRC-NEXT: v_mov_b32_e32 v4, s20 +; HEURRC-NEXT: v_mov_b32_e32 v5, s21 +; HEURRC-NEXT: v_mov_b32_e32 v6, s22 +; HEURRC-NEXT: v_mov_b32_e32 v7, s23 +; HEURRC-NEXT: v_mov_b32_e32 v8, s24 +; HEURRC-NEXT: v_mov_b32_e32 v9, s25 +; HEURRC-NEXT: v_mov_b32_e32 v10, s26 +; HEURRC-NEXT: v_mov_b32_e32 v11, s27 +; HEURRC-NEXT: v_mov_b32_e32 v12, s28 +; HEURRC-NEXT: v_mov_b32_e32 v13, s29 +; HEURRC-NEXT: v_mov_b32_e32 v14, s30 +; HEURRC-NEXT: v_mov_b32_e32 v15, s31 +; HEURRC-NEXT: v_mov_b32_e32 v16, s0 +; HEURRC-NEXT: v_mov_b32_e32 v17, s1 +; HEURRC-NEXT: v_mov_b32_e32 v18, s2 +; HEURRC-NEXT: v_mov_b32_e32 v19, s3 +; HEURRC-NEXT: v_mov_b32_e32 v20, s4 +; HEURRC-NEXT: v_mov_b32_e32 v21, s5 +; HEURRC-NEXT: v_mov_b32_e32 v22, s6 +; HEURRC-NEXT: v_mov_b32_e32 v23, s7 +; HEURRC-NEXT: v_mov_b32_e32 v24, s8 +; HEURRC-NEXT: v_mov_b32_e32 v25, s9 +; HEURRC-NEXT: v_mov_b32_e32 v26, s10 +; HEURRC-NEXT: v_mov_b32_e32 v27, s11 +; HEURRC-NEXT: v_mov_b32_e32 v28, s12 +; HEURRC-NEXT: v_mov_b32_e32 v29, s13 +; HEURRC-NEXT: v_mov_b32_e32 v30, s14 +; HEURRC-NEXT: v_mov_b32_e32 v31, s15 +; HEURRC-NEXT: s_nop 1 +; HEURRC-NEXT: v_mfma_f32_32x32x1_2b_f32 v[0:31], v33, v34, v[0:31] cbsz:1 abid:2 blgp:3 +; HEURRC-NEXT: s_nop 15 +; HEURRC-NEXT: s_nop 1 +; HEURRC-NEXT: global_store_dwordx4 v32, v[24:27], s[34:35] offset:96 +; HEURRC-NEXT: global_store_dwordx4 v32, v[28:31], s[34:35] offset:112 +; HEURRC-NEXT: global_store_dwordx4 v32, v[16:19], s[34:35] offset:64 +; HEURRC-NEXT: global_store_dwordx4 v32, v[20:23], s[34:35] offset:80 +; HEURRC-NEXT: global_store_dwordx4 v32, v[8:11], s[34:35] offset:32 +; HEURRC-NEXT: global_store_dwordx4 v32, v[12:15], s[34:35] offset:48 +; HEURRC-NEXT: global_store_dwordx4 v32, v[0:3], s[34:35] +; HEURRC-NEXT: global_store_dwordx4 v32, v[4:7], s[34:35] offset:16 +; HEURRC-NEXT: s_endpgm +; +; VGPRRC-LABEL: not_enough_agprs: +; VGPRRC: ; %bb.0: ; %bb +; VGPRRC-NEXT: s_load_dwordx2 s[34:35], s[4:5], 0x24 +; VGPRRC-NEXT: v_mov_b32_e32 v33, 1.0 +; VGPRRC-NEXT: v_mov_b32_e32 v34, 2.0 +; VGPRRC-NEXT: v_mov_b32_e32 v32, 0 +; VGPRRC-NEXT: s_waitcnt lgkmcnt(0) +; VGPRRC-NEXT: s_load_dwordx16 s[16:31], s[34:35], 0x0 +; VGPRRC-NEXT: s_load_dwordx16 s[0:15], s[34:35], 0x40 +; VGPRRC-NEXT: s_waitcnt lgkmcnt(0) +; VGPRRC-NEXT: v_mov_b32_e32 v0, s16 +; VGPRRC-NEXT: v_mov_b32_e32 v1, s17 +; VGPRRC-NEXT: v_mov_b32_e32 v2, s18 +; VGPRRC-NEXT: v_mov_b32_e32 v3, s19 +; VGPRRC-NEXT: v_mov_b32_e32 v4, s20 +; VGPRRC-NEXT: v_mov_b32_e32 v5, s21 +; VGPRRC-NEXT: v_mov_b32_e32 v6, s22 +; VGPRRC-NEXT: v_mov_b32_e32 v7, s23 +; VGPRRC-NEXT: v_mov_b32_e32 v8, s24 +; VGPRRC-NEXT: v_mov_b32_e32 v9, s25 +; VGPRRC-NEXT: v_mov_b32_e32 v10, s26 +; VGPRRC-NEXT: v_mov_b32_e32 v11, s27 +; VGPRRC-NEXT: v_mov_b32_e32 v12, s28 +; VGPRRC-NEXT: v_mov_b32_e32 v13, s29 +; VGPRRC-NEXT: v_mov_b32_e32 v14, s30 +; VGPRRC-NEXT: v_mov_b32_e32 v15, s31 +; VGPRRC-NEXT: v_mov_b32_e32 v16, s0 +; VGPRRC-NEXT: v_mov_b32_e32 v17, s1 +; VGPRRC-NEXT: v_mov_b32_e32 v18, s2 +; VGPRRC-NEXT: v_mov_b32_e32 v19, s3 +; VGPRRC-NEXT: v_mov_b32_e32 v20, s4 +; VGPRRC-NEXT: v_mov_b32_e32 v21, s5 +; VGPRRC-NEXT: v_mov_b32_e32 v22, s6 +; VGPRRC-NEXT: v_mov_b32_e32 v23, s7 +; VGPRRC-NEXT: v_mov_b32_e32 v24, s8 +; VGPRRC-NEXT: v_mov_b32_e32 v25, s9 +; VGPRRC-NEXT: v_mov_b32_e32 v26, s10 +; VGPRRC-NEXT: v_mov_b32_e32 v27, s11 +; VGPRRC-NEXT: v_mov_b32_e32 v28, s12 +; VGPRRC-NEXT: v_mov_b32_e32 v29, s13 +; VGPRRC-NEXT: v_mov_b32_e32 v30, s14 +; VGPRRC-NEXT: v_mov_b32_e32 v31, s15 +; VGPRRC-NEXT: s_nop 1 +; VGPRRC-NEXT: v_mfma_f32_32x32x1_2b_f32 v[0:31], v33, v34, v[0:31] cbsz:1 abid:2 blgp:3 +; VGPRRC-NEXT: s_nop 15 +; VGPRRC-NEXT: s_nop 1 +; VGPRRC-NEXT: global_store_dwordx4 v32, v[24:27], s[34:35] offset:96 +; VGPRRC-NEXT: global_store_dwordx4 v32, v[28:31], s[34:35] offset:112 +; VGPRRC-NEXT: global_store_dwordx4 v32, v[16:19], s[34:35] offset:64 +; VGPRRC-NEXT: global_store_dwordx4 v32, v[20:23], s[34:35] offset:80 +; VGPRRC-NEXT: global_store_dwordx4 v32, v[8:11], s[34:35] offset:32 +; VGPRRC-NEXT: global_store_dwordx4 v32, v[12:15], s[34:35] offset:48 +; VGPRRC-NEXT: global_store_dwordx4 v32, v[0:3], s[34:35] +; VGPRRC-NEXT: global_store_dwordx4 v32, v[4:7], s[34:35] offset:16 +; VGPRRC-NEXT: s_endpgm +bb: + %in.1 = load <32 x float>, ptr addrspace(1) %arg, align 128 + %mai.1 = call <32 x float> @llvm.amdgcn.mfma.f32.32x32x1f32(float 1.000000e+00, float 2.000000e+00, <32 x float> %in.1, i32 1, i32 2, i32 3) + store <32 x float> %mai.1, ptr addrspace(1) %arg, align 128 + ret void +} + +define <16 x float> @mfma_scale_respect_flag(<8 x i32> %arg0, <8 x i32> %arg1, <16 x float> %arg2, i32 %scale0, i32 %scale1) #2 { +; HEURRC-LABEL: mfma_scale_respect_flag: +; HEURRC: ; %bb.0: +; HEURRC-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) +; HEURRC-NEXT: scratch_load_dword a15, off, s32 +; HEURRC-NEXT: scratch_load_dword v31, off, s32 offset:8 +; HEURRC-NEXT: scratch_load_dword v32, off, s32 offset:4 +; HEURRC-NEXT: v_accvgpr_write_b32 a0, v16 +; HEURRC-NEXT: v_accvgpr_write_b32 a1, v17 +; HEURRC-NEXT: v_accvgpr_write_b32 a2, v18 +; HEURRC-NEXT: v_accvgpr_write_b32 a3, v19 +; HEURRC-NEXT: v_accvgpr_write_b32 a4, v20 +; HEURRC-NEXT: v_accvgpr_write_b32 a5, v21 +; HEURRC-NEXT: v_accvgpr_write_b32 a6, v22 +; HEURRC-NEXT: v_accvgpr_write_b32 a7, v23 +; HEURRC-NEXT: v_accvgpr_write_b32 a8, v24 +; HEURRC-NEXT: v_accvgpr_write_b32 a9, v25 +; HEURRC-NEXT: v_accvgpr_write_b32 a10, v26 +; HEURRC-NEXT: v_accvgpr_write_b32 a11, v27 +; HEURRC-NEXT: v_accvgpr_write_b32 a12, v28 +; HEURRC-NEXT: v_accvgpr_write_b32 a13, v29 +; HEURRC-NEXT: v_accvgpr_write_b32 a14, v30 +; HEURRC-NEXT: s_waitcnt vmcnt(0) +; HEURRC-NEXT: s_nop 0 +; HEURRC-NEXT: v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], v[0:7], v[8:15], a[0:15], v32, v31 op_sel_hi:[0,0,0] +; HEURRC-NEXT: s_nop 15 +; HEURRC-NEXT: s_nop 3 +; HEURRC-NEXT: v_accvgpr_read_b32 v0, a0 +; HEURRC-NEXT: v_accvgpr_read_b32 v1, a1 +; HEURRC-NEXT: v_accvgpr_read_b32 v2, a2 +; HEURRC-NEXT: v_accvgpr_read_b32 v3, a3 +; HEURRC-NEXT: v_accvgpr_read_b32 v4, a4 +; HEURRC-NEXT: v_accvgpr_read_b32 v5, a5 +; HEURRC-NEXT: v_accvgpr_read_b32 v6, a6 +; HEURRC-NEXT: v_accvgpr_read_b32 v7, a7 +; HEURRC-NEXT: v_accvgpr_read_b32 v8, a8 +; HEURRC-NEXT: v_accvgpr_read_b32 v9, a9 +; HEURRC-NEXT: v_accvgpr_read_b32 v10, a10 +; HEURRC-NEXT: v_accvgpr_read_b32 v11, a11 +; HEURRC-NEXT: v_accvgpr_read_b32 v12, a12 +; HEURRC-NEXT: v_accvgpr_read_b32 v13, a13 +; HEURRC-NEXT: v_accvgpr_read_b32 v14, a14 +; HEURRC-NEXT: v_accvgpr_read_b32 v15, a15 +; HEURRC-NEXT: s_setpc_b64 s[30:31] +; +; VGPRRC-LABEL: mfma_scale_respect_flag: +; VGPRRC: ; %bb.0: +; VGPRRC-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0) +; VGPRRC-NEXT: scratch_load_dword v31, off, s32 +; VGPRRC-NEXT: scratch_load_dword v32, off, s32 offset:8 +; VGPRRC-NEXT: scratch_load_dword v33, off, s32 offset:4 +; VGPRRC-NEXT: s_waitcnt vmcnt(0) +; VGPRRC-NEXT: v_mfma_scale_f32_32x32x64_f8f6f4 v[16:31], v[0:7], v[8:15], v[16:31], v33, v32 op_sel_hi:[0,0,0] +; VGPRRC-NEXT: s_nop 15 +; VGPRRC-NEXT: s_nop 3 +; VGPRRC-NEXT: v_mov_b32_e32 v0, v16 +; VGPRRC-NEXT: v_mov_b32_e32 v1, v17 +; VGPRRC-NEXT: v_mov_b32_e32 v2, v18 +; VGPRRC-NEXT: v_mov_b32_e32 v3, v19 +; VGPRRC-NEXT: v_mov_b32_e32 v4, v20 +; VGPRRC-NEXT: v_mov_b32_e32 v5, v21 +; VGPRRC-NEXT: v_mov_b32_e32 v6, v22 +; VGPRRC-NEXT: v_mov_b32_e32 v7, v23 +; VGPRRC-NEXT: v_mov_b32_e32 v8, v24 +; VGPRRC-NEXT: v_mov_b32_e32 v9, v25 +; VGPRRC-NEXT: v_mov_b32_e32 v10, v26 +; VGPRRC-NEXT: v_mov_b32_e32 v11, v27 +; VGPRRC-NEXT: v_mov_b32_e32 v12, v28 +; VGPRRC-NEXT: v_mov_b32_e32 v13, v29 +; VGPRRC-NEXT: v_mov_b32_e32 v14, v30 +; VGPRRC-NEXT: v_mov_b32_e32 v15, v31 +; VGPRRC-NEXT: s_setpc_b64 s[30:31] + %result = call <16 x float> @llvm.amdgcn.mfma.scale.f32.32x32x64.f8f6f4.v8i32.v8i32(<8 x i32> %arg0, <8 x i32> %arg1, <16 x float> %arg2, + i32 0, ; cbsz + i32 0, ; blgp + i32 0, i32 %scale0, i32 0, i32 %scale1) + ret <16 x float> %result +} + attributes #0 = { "amdgpu-agpr-alloc"="32,256" } attributes #1 = { "amdgpu-agpr-alloc"="0,0" } +attributes #2 = { nounwind "amdgpu-agpr-alloc"="20" }