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<SIMachineFunctionInfo>();
       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<SIMachineFunctionInfo>();
       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 908d856d386f5..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<bool> MFMAVGPRForm(
-    "amdgpu-mfma-vgpr-form", cl::Hidden,
+static cl::opt<bool, true> 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<const GCNTargetMachine &>(TLI->getTargetMachine());
 }
 
+bool SIMachineFunctionInfo::MFMAVGPRForm = false;
+
 SIMachineFunctionInfo::SIMachineFunctionInfo(const Function &F,
                                              const GCNSubtarget *STI)
     : AMDGPUMachineFunction(F, *STI), Mode(F, *STI), GWSResourcePSV(getTM(STI)),
@@ -81,14 +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 (MFMAVGPRForm ||
-        (ST.getMaxNumVGPRs(F) <= ST.getAddressableNumArchVGPRs() &&
-         !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<MCPhysReg, 32> 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<VOPProfile P, VOP3Features Features = VOP3_REGULAR,
 class VOP3P_Mix_Profile_t16<VOPProfile P, VOP3Features Features = VOP3_REGULAR>
      :  VOP3P_Mix_Profile<P, Features, 0> {
   let IsTrue16 = 1;
-  let IsRealTrue16 = 1; 
+  let IsRealTrue16 = 1;
   let DstRC64 = getVALUDstForVT<P.DstVT, 1 /*IsTrue16*/, 1 /*IsVOP3Encoding*/>.ret;
 }
 
@@ -950,7 +950,7 @@ class MFMA_F8F6F4_WithSizeTable_Helper<VOP3_Pseudo  ps, string F8F8Op> :
 }
 
 // Currently assumes scaled instructions never have abid
-class MAIFrag<SDPatternOperator Op, code pred, bit HasAbid = true, bit Scaled = false> : PatFrag <
+class MAIFrag<SDPatternOperator Op, bit HasAbid = true, bit Scaled = false> : 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<SDPatternOperator Op, code pred, bit HasAbid = true, bit Scaled =
                     (ops node:$blgp))),
   !if(Scaled, (Op $src0, $src1, $src2, $cbsz, $blgp, $src0_modifiers, $scale_src0, $src1_modifiers, $scale_src1),
               !if(HasAbid, (Op $src0, $src1, $src2, $cbsz, $abid, $blgp),
-                           (Op $src0, $src1, $src2, $cbsz, $blgp))),
-  pred
->;
-
-defvar MayNeedAGPRs = [{
-  return MF->getInfo<SIMachineFunctionInfo>()->mayNeedAGPRs();
-}];
-
-defvar MayNeedAGPRs_gisel = [{
-  return MF.getInfo<SIMachineFunctionInfo>()->mayNeedAGPRs();
-}];
+                           (Op $src0, $src1, $src2, $cbsz, $blgp)))>;
 
-defvar MayNotNeedAGPRs = [{
-  return !MF->getInfo<SIMachineFunctionInfo>()->mayNeedAGPRs();
-}];
+class CanUseAGPR_MAI<ValueType vt> {
+  code PredicateCode = [{
+    return !Subtarget->hasGFX90AInsts() ||
+      (!SIMachineFunctionInfo::MFMAVGPRForm &&
+       MF->getInfo<SIMachineFunctionInfo>()->getMinNumAGPRs() >=
+  }] # !srl(vt.Size, 5) # ");";
 
-defvar MayNotNeedAGPRs_gisel = [{
-  return !MF.getInfo<SIMachineFunctionInfo>()->mayNeedAGPRs();
-}];
+  code GISelPredicateCode = [{
+    return !Subtarget->hasGFX90AInsts() ||
+      (!SIMachineFunctionInfo::MFMAVGPRForm &&
+       MF.getInfo<SIMachineFunctionInfo>()->getMinNumAGPRs() >=
+  }] # !srl(vt.Size, 5) # ");";
+}
 
-class AgprMAIFrag<SDPatternOperator Op, bit HasAbid = true,
+class AgprMAIFrag<SDPatternOperator Op, ValueType vt, bit HasAbid = true,
                   bit Scaled = false> :
-  MAIFrag<Op, MayNeedAGPRs, HasAbid, Scaled> {
-  let GISelPredicateCode = MayNeedAGPRs_gisel;
-}
+  MAIFrag<Op, HasAbid, Scaled>,
+  CanUseAGPR_MAI<vt>;
 
 class VgprMAIFrag<SDPatternOperator Op, bit HasAbid = true,
-                   bit Scaled = false> :
-  MAIFrag<Op, MayNotNeedAGPRs, HasAbid, Scaled> {
-  let GISelPredicateCode = MayNotNeedAGPRs_gisel;
-}
+                  bit Scaled = false> :
+  MAIFrag<Op, HasAbid, Scaled>;
 
 let isAsCheapAsAMove = 1, isReMaterializable = 1 in {
   defm V_ACCVGPR_READ_B32  : VOP3Inst<"v_accvgpr_read_b32",  VOPProfileAccRead>;
@@ -1037,16 +1030,19 @@ multiclass MAIInst<string OpName, string P, SDPatternOperator node = null_frag,
                    bit HasAbid = true,
                    bit Scaled = false> {
   defvar NoDstOverlap = !cast<VOPProfileMAI>("VOPProfileMAI_" # P).NoDstOverlap;
+  defvar ProfileAGPR = !cast<VOPProfileMAI>("VOPProfileMAI_" # P);
+  defvar ProfileVGPR = !cast<VOPProfileMAI>("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<OpName, !cast<VOPProfileMAI>("VOPProfileMAI_" # P),
-                         !if(!or(NoDstOverlap, !eq(node, null_frag)), null_frag, AgprMAIFrag<node, HasAbid, Scaled>), Scaled>,
+      def _e64 : MAIInst<OpName, ProfileAGPR,
+                         !if(!or(NoDstOverlap, !eq(node, null_frag)), null_frag, AgprMAIFrag<node, ProfileAGPR.DstVT, HasAbid, Scaled>), Scaled>,
                  MFMATable<0, "AGPR", NAME # "_e64">;
 
       let OtherPredicates = [isGFX90APlus], Mnemonic = OpName in
-      def _vgprcd_e64 : MAIInst<OpName # "_vgprcd", !cast<VOPProfileMAI>("VOPProfileMAI_" # P # "_VCD"),
+      def _vgprcd_e64 : MAIInst<OpName # "_vgprcd", ProfileVGPR,
                                 !if(!or(NoDstOverlap, !eq(node, null_frag)), null_frag, VgprMAIFrag<node, HasAbid, Scaled>), Scaled>,
                         MFMATable<0, "VGPR", NAME # "_vgprcd_e64", NAME # "_e64">;
     }
@@ -1055,12 +1051,12 @@ multiclass MAIInst<string OpName, string P, SDPatternOperator node = null_frag,
       let Constraints = !if(NoDstOverlap, "$vdst = $src2", ""),
           isConvertibleToThreeAddress = NoDstOverlap,
           Mnemonic = OpName in {
-        def "_mac_e64" : MAIInst<OpName # "_mac", !cast<VOPProfileMAI>("VOPProfileMAI_" # P),
-                                 !if(!eq(node, null_frag), null_frag, AgprMAIFrag<node, HasAbid, Scaled>), Scaled>,
+        def "_mac_e64" : MAIInst<OpName # "_mac", ProfileAGPR,
+                                 !if(!eq(node, null_frag), null_frag, AgprMAIFrag<node, ProfileAGPR.DstVT, HasAbid, Scaled>), Scaled>,
                          MFMATable<1, "AGPR", NAME # "_e64", NAME # "_mac_e64">;
 
         let OtherPredicates = [isGFX90APlus] in
-        def _mac_vgprcd_e64 : MAIInst<OpName # "_mac_vgprcd", !cast<VOPProfileMAI>("VOPProfileMAI_" # P # "_VCD"),
+        def _mac_vgprcd_e64 : MAIInst<OpName # "_mac_vgprcd", ProfileVGPR,
                                       !if(!eq(node, null_frag), null_frag, VgprMAIFrag<node, HasAbid, Scaled>), Scaled>,
                               MFMATable<1, "VGPR", NAME # "_vgprcd_e64", NAME # "_mac_e64">;
       }
@@ -1074,11 +1070,11 @@ multiclass ScaledMAIInst_mc<string OpName, string UnscaledOpName_, SDPatternOper
   defvar UnscaledOpName = UnscaledOpName_#VariantSuffix;
 
   defvar HasAbid = false;
-
-  defvar NoDstOverlap = !cast<VOPProfileMAI>(!cast<MAIInst>(UnscaledOpName#"_e64").Pfl).NoDstOverlap;
+  defvar Profile = !cast<VOPProfileMAI>(!cast<MAIInst>(UnscaledOpName#"_e64").Pfl);
+  defvar NoDstOverlap = Profile.NoDstOverlap;
 
   def _e64 : ScaledMAIInst<OpName,
-        !cast<MAIInst>(UnscaledOpName#"_e64"), !if(NoDstOverlap, null_frag, AgprMAIFrag<node, HasAbid, true>)>,
+        !cast<MAIInst>(UnscaledOpName#"_e64"), !if(NoDstOverlap, null_frag, AgprMAIFrag<node, Profile.DstVT, HasAbid, true>)>,
       MFMATable<0, "AGPR", NAME # "_e64">;
 
   def _vgprcd_e64 : ScaledMAIInst<OpName # "_vgprcd",
@@ -1090,7 +1086,7 @@ multiclass ScaledMAIInst_mc<string OpName, string UnscaledOpName_, SDPatternOper
        isConvertibleToThreeAddress = NoDstOverlap,
        Mnemonic = UnscaledOpName_ in {
      def _mac_e64 : ScaledMAIInst<OpName # "_mac",
-          !cast<MAIInst>(UnscaledOpName # "_mac_e64"), AgprMAIFrag<node, HasAbid, true>>,
+          !cast<MAIInst>(UnscaledOpName # "_mac_e64"), AgprMAIFrag<node, Profile.DstVT, HasAbid, true>>,
         MFMATable<1, "AGPR", NAME # "_e64">;
 
      def _mac_vgprcd_e64 : ScaledMAIInst<OpName # " _mac_vgprcd",
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.form.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.form.ll
index 87a7c2ef6c95c..cc4cc8efd66f4 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.form.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.mfma.form.ll
@@ -72,5 +72,206 @@ define <4 x float> @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" }
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" }