[AMDGPU] gfx11 VINTERP intrinsics and ISel support

[Joe_Nash]

Summary

Depends on D127664

[rampitec]

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

1478

Are you going to submit these builtins separately?

[Joe_Nash]

removed builtins

[rampitec]

LGTM

[foad]

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

1478

The __int_* prefix doesn't make much sense. I would suggest either using the tablegen name (int_amdgcn_interp_inreg_p10) or preferably the LLVM IR name (llvm.amdgcn.interp.inreg.p10).

[Joe_Nash]

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

1478

Are you going to submit these builtins separately?

I will remove the builtin and it can be added later if needed.

1478

Done, see
75378d432fda5408b7210fd3627db884561db650

[kosarev]

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

1494

Wouldn't it be more natural to declare p and p0 to be llvm_v2f16_ty?

[foad]

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

1494

Yes, I think so, but it would require coordinated changes in Mesa and LLPC, and we should probably also provide a v2f16 version of llvm.amdgcn.lds.param.load which currently always returns a float.

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

 // so it behaves like IntrNoMem.
 def int_amdgcn_lds_param_load :
   GCCBuiltin<"__builtin_amdgcn_lds_param_load">,
   Intrinsic<[llvm_float_ty],
             [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
             [IntrNoMem, IntrSpeculatable, IntrWillReturn,
              ImmArg<ArgIndex<0>>, ImmArg<ArgIndex<1>>]>;
 
+// __int_amdgcn_interp_inreg_p10 <p>, <i>, <p0>
+def int_amdgcn_interp_inreg_p10 :
+  Intrinsic<[llvm_float_ty],
+            [llvm_float_ty, llvm_float_ty, llvm_float_ty],
+            [IntrNoMem, IntrSpeculatable, IntrWillReturn]>;
+
+// __int_amdgcn_interp_inreg_p2 <p>, <j>, <tmp>
+def int_amdgcn_interp_inreg_p2 :
+  Intrinsic<[llvm_float_ty],
+            [llvm_float_ty, llvm_float_ty, llvm_float_ty],
+            [IntrNoMem, IntrSpeculatable, IntrWillReturn]>;
+
+// __int_amdgcn_interp_inreg_p10_f16 <p>, <i>, <p0>, <high>
+// high selects whether high or low 16-bits are used for p and p0 operands
+def int_amdgcn_interp_inreg_p10_f16:
+  Intrinsic<[llvm_float_ty],
+            [llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_i1_ty],
+            [IntrNoMem, IntrSpeculatable, IntrWillReturn,
+             ImmArg<ArgIndex<3>>]>;
+
+// __int_amdgcn_interp_inreg_p2_f16 <p>, <j>, <tmp>, <high>
+// high selects whether high or low 16-bits are used for p operand
+def int_amdgcn_interp_inreg_p2_f16 :
+  Intrinsic<[llvm_half_ty],
+            [llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_i1_ty],
+            [IntrNoMem, IntrSpeculatable, IntrWillReturn,
+             ImmArg<ArgIndex<3>>]>;
+
 // Deprecated: use llvm.amdgcn.live.mask instead.
 def int_amdgcn_ps_live : Intrinsic <
   [llvm_i1_ty],
   [],
   [IntrNoMem, IntrWillReturn]>;
 
 // Query currently live lanes.
 // Returns true if lane is live (and not a helper lane).

llvm/lib/Target/AMDGPU/AMDGPUGISel.td

 def gi_vop3pmodsdot :
     GIComplexOperandMatcher<s32, "selectVOP3PModsDOT">,
     GIComplexPatternEquiv<VOP3PModsDOT>;
 
 def gi_vop3opselmods :
     GIComplexOperandMatcher<s32, "selectVOP3OpSelMods">,
     GIComplexPatternEquiv<VOP3OpSelMods>;
 
+def gi_vinterpmods :
+    GIComplexOperandMatcher<s32, "selectVINTERPMods">,
+    GIComplexPatternEquiv<VINTERPMods>;
+
+def gi_vinterpmods_hi :
+    GIComplexOperandMatcher<s32, "selectVINTERPModsHi">,
+    GIComplexPatternEquiv<VINTERPModsHi>;
+
 // FIXME: Why do we have both VOP3OpSel and VOP3OpSelMods?
 def gi_vop3opsel :
     GIComplexOperandMatcher<s32, "selectVOP3OpSelMods">,
     GIComplexPatternEquiv<VOP3OpSel>;
 
 def gi_smrd_imm :
     GIComplexOperandMatcher<s64, "selectSmrdImm">,
     GIComplexPatternEquiv<SMRDImm>;

llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.h

private:
   bool SelectVOP3NoMods(SDValue In, SDValue &Src) const;
   bool SelectVOP3Mods0(SDValue In, SDValue &Src, SDValue &SrcMods,
                        SDValue &Clamp, SDValue &Omod) const;
   bool SelectVOP3BMods0(SDValue In, SDValue &Src, SDValue &SrcMods,
                         SDValue &Clamp, SDValue &Omod) const;
   bool SelectVOP3NoMods0(SDValue In, SDValue &Src, SDValue &SrcMods,
                          SDValue &Clamp, SDValue &Omod) const;
 
+  bool SelectVINTERPModsImpl(SDValue In, SDValue &Src, SDValue &SrcMods,
+                             bool OpSel) const;
+  bool SelectVINTERPMods(SDValue In, SDValue &Src, SDValue &SrcMods) const;
+  bool SelectVINTERPModsHi(SDValue In, SDValue &Src, SDValue &SrcMods) const;
+
   bool SelectVOP3OMods(SDValue In, SDValue &Src, SDValue &Clamp,
                        SDValue &Omod) const;
 
   bool SelectVOP3PMods(SDValue In, SDValue &Src, SDValue &SrcMods,
                        bool IsDOT = false) const;
   bool SelectVOP3PModsDOT(SDValue In, SDValue &Src, SDValue &SrcMods) const;
 
   bool SelectVOP3OpSel(SDValue In, SDValue &Src, SDValue &SrcMods) const;

llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp

 //===-- AMDGPUISelDAGToDAG.cpp - A dag to dag inst selector for AMDGPU ----===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //==-----------------------------------------------------------------------===//
 //
 /// \file
 /// Defines an instruction selector for the AMDGPU target.
 //
 //===----------------------------------------------------------------------===//
 
 #include "AMDGPUISelDAGToDAG.h"
 #include "AMDGPU.h"
+#include "AMDGPUSubtarget.h"
 #include "AMDGPUTargetMachine.h"
+#include "MCTargetDesc/AMDGPUMCTargetDesc.h"
 #include "MCTargetDesc/R600MCTargetDesc.h"
 #include "R600RegisterInfo.h"
 #include "SIMachineFunctionInfo.h"
 #include "llvm/Analysis/LegacyDivergenceAnalysis.h"
 #include "llvm/Analysis/ValueTracking.h"
 #include "llvm/CodeGen/FunctionLoweringInfo.h"
 #include "llvm/CodeGen/SelectionDAG.h"
 #include "llvm/CodeGen/SelectionDAGISel.h"

 bool AMDGPUDAGToDAGISel::SelectVOP3NoMods(SDValue In, SDValue &Src) const {
   if (In.getOpcode() == ISD::FABS || In.getOpcode() == ISD::FNEG)
     return false;
 
   Src = In;
   return true;
 }
 
+bool AMDGPUDAGToDAGISel::SelectVINTERPModsImpl(SDValue In, SDValue &Src,
+                                               SDValue &SrcMods,
+                                               bool OpSel) const {
+  unsigned Mods;
+  if (SelectVOP3ModsImpl(In, Src, Mods, /* AllowAbs */ false)) {
+    if (OpSel)
+      Mods |= SISrcMods::OP_SEL_0;
+    SrcMods = CurDAG->getTargetConstant(Mods, SDLoc(In), MVT::i32);
+    return true;
+  }
+
+  return false;
+}
+
+bool AMDGPUDAGToDAGISel::SelectVINTERPMods(SDValue In, SDValue &Src,
+                                           SDValue &SrcMods) const {
+  return SelectVINTERPModsImpl(In, Src, SrcMods, /* OpSel */ false);
+}
+
+bool AMDGPUDAGToDAGISel::SelectVINTERPModsHi(SDValue In, SDValue &Src,
+                                             SDValue &SrcMods) const {
+  return SelectVINTERPModsImpl(In, Src, SrcMods, /* OpSel */ true);
+}
+
 bool AMDGPUDAGToDAGISel::SelectVOP3Mods0(SDValue In, SDValue &Src,
                                          SDValue &SrcMods, SDValue &Clamp,
                                          SDValue &Omod) const {
   SDLoc DL(In);
   Clamp = CurDAG->getTargetConstant(0, DL, MVT::i1);
   Omod = CurDAG->getTargetConstant(0, DL, MVT::i1);
 
   return SelectVOP3Mods(In, Src, SrcMods);

llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.h

private:
   bool selectGlobalAtomicFadd(MachineInstr &I, MachineOperand &AddrOp,
                               MachineOperand &DataOp) const;
   bool selectBufferLoadLds(MachineInstr &MI) const;
   bool selectGlobalLoadLds(MachineInstr &MI) const;
   bool selectBVHIntrinsic(MachineInstr &I) const;
   bool selectSMFMACIntrin(MachineInstr &I) const;
   bool selectWaveAddress(MachineInstr &I) const;
 
-  std::pair<Register, unsigned> selectVOP3ModsImpl(MachineOperand &Root,
-                                                   bool AllowAbs = true) const;
+  std::pair<Register, unsigned>
+  selectVOP3ModsImpl(MachineOperand &Root, bool AllowAbs = true,
+                     bool OpSel = false, bool ForceVGPR = false) const;
 
   InstructionSelector::ComplexRendererFns
   selectVCSRC(MachineOperand &Root) const;
 
   InstructionSelector::ComplexRendererFns
   selectVSRC0(MachineOperand &Root) const;
 
   InstructionSelector::ComplexRendererFns
private:
 
   InstructionSelector::ComplexRendererFns
   selectVOP3PModsDOT(MachineOperand &Root) const;
 
   InstructionSelector::ComplexRendererFns
   selectVOP3OpSelMods(MachineOperand &Root) const;
 
   InstructionSelector::ComplexRendererFns
+  selectVINTERPMods(MachineOperand &Root) const;
+  InstructionSelector::ComplexRendererFns
+  selectVINTERPModsHi(MachineOperand &Root) const;
+
+  InstructionSelector::ComplexRendererFns
   selectSmrdImm(MachineOperand &Root) const;
   InstructionSelector::ComplexRendererFns
   selectSmrdImm32(MachineOperand &Root) const;
   InstructionSelector::ComplexRendererFns
   selectSmrdSgpr(MachineOperand &Root) const;
 
   std::pair<Register, int> selectFlatOffsetImpl(MachineOperand &Root,
                                                 uint64_t FlatVariant) const;

llvm/lib/Target/AMDGPU/AMDGPUInstructionSelector.cpp

 InstructionSelector::ComplexRendererFns
 AMDGPUInstructionSelector::selectVCSRC(MachineOperand &Root) const {
   return {{
       [=](MachineInstrBuilder &MIB) { MIB.add(Root); }
   }};
 
 }
 
-std::pair<Register, unsigned>
-AMDGPUInstructionSelector::selectVOP3ModsImpl(MachineOperand &Root,
-                                              bool AllowAbs) const {
+std::pair<Register, unsigned> AMDGPUInstructionSelector::selectVOP3ModsImpl(
+    MachineOperand &Root, bool AllowAbs, bool OpSel, bool ForceVGPR) const {
   Register Src = Root.getReg();
   Register OrigSrc = Src;
   unsigned Mods = 0;
   MachineInstr *MI = getDefIgnoringCopies(Src, *MRI);
 
   if (MI && MI->getOpcode() == AMDGPU::G_FNEG) {
     Src = MI->getOperand(1).getReg();
     Mods |= SISrcMods::NEG;
     MI = getDefIgnoringCopies(Src, *MRI);
   }
 
   if (AllowAbs && MI && MI->getOpcode() == AMDGPU::G_FABS) {
     Src = MI->getOperand(1).getReg();
     Mods |= SISrcMods::ABS;
   }
 
-  if (Mods != 0 &&
+  if (OpSel)
+    Mods |= SISrcMods::OP_SEL_0;
+
+  if ((Mods != 0 || ForceVGPR) &&
       RBI.getRegBank(Src, *MRI, TRI)->getID() != AMDGPU::VGPRRegBankID) {
     MachineInstr *UseMI = Root.getParent();
 
     // If we looked through copies to find source modifiers on an SGPR operand,
     // we now have an SGPR register source. To avoid potentially violating the
     // constant bus restriction, we need to insert a copy to a VGPR.
     Register VGPRSrc = MRI->cloneVirtualRegister(OrigSrc);
     BuildMI(*UseMI->getParent(), UseMI, UseMI->getDebugLoc(),
AMDGPUInstructionSelector::selectVOP3OpSelMods(MachineOperand &Root) const {
   // FIXME: Handle op_sel
   return {{
       [=](MachineInstrBuilder &MIB) { MIB.addReg(Root.getReg()); },
       [=](MachineInstrBuilder &MIB) { MIB.addImm(0); } // src_mods
   }};
 }
 
 InstructionSelector::ComplexRendererFns
+AMDGPUInstructionSelector::selectVINTERPMods(MachineOperand &Root) const {
+  Register Src;
+  unsigned Mods;
+  std::tie(Src, Mods) = selectVOP3ModsImpl(Root,
+                                           /* AllowAbs */ false,
+                                           /* OpSel */ false,
+                                           /* ForceVGPR */ true);
+
+  return {{
+      [=](MachineInstrBuilder &MIB) { MIB.addReg(Src); },
+      [=](MachineInstrBuilder &MIB) { MIB.addImm(Mods); }, // src0_mods
+  }};
+}
+
+InstructionSelector::ComplexRendererFns
+AMDGPUInstructionSelector::selectVINTERPModsHi(MachineOperand &Root) const {
+  Register Src;
+  unsigned Mods;
+  std::tie(Src, Mods) = selectVOP3ModsImpl(Root,
+                                           /* AllowAbs */ false,
+                                           /* OpSel */ true,
+                                           /* ForceVGPR */ true);
+
+  return {{
+      [=](MachineInstrBuilder &MIB) { MIB.addReg(Src); },
+      [=](MachineInstrBuilder &MIB) { MIB.addImm(Mods); }, // src0_mods
+  }};
+}
+
+InstructionSelector::ComplexRendererFns
 AMDGPUInstructionSelector::selectSmrdImm(MachineOperand &Root) const {
   SmallVector<GEPInfo, 4> AddrInfo;
   getAddrModeInfo(*Root.getParent(), *MRI, AddrInfo);
 
   if (AddrInfo.empty() || AddrInfo[0].SgprParts.size() != 1)
     return None;
 
   const GEPInfo &GEPInfo = AddrInfo[0];

llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp

case AMDGPU::G_INTRINSIC: {
     case Intrinsic::amdgcn_lds_param_load: {
       applyDefaultMapping(OpdMapper);
 
       // Readlane for m0 value, which is always the last operand.
       // FIXME: Should this be a waterfall loop instead?
       constrainOpWithReadfirstlane(MI, MRI, MI.getNumOperands() - 1); // Index
       return;
     }
+    case Intrinsic::amdgcn_interp_inreg_p10:
+    case Intrinsic::amdgcn_interp_inreg_p2:
+    case Intrinsic::amdgcn_interp_inreg_p10_f16:
+    case Intrinsic::amdgcn_interp_inreg_p2_f16:
+      applyDefaultMapping(OpdMapper);
+      return;
     case Intrinsic::amdgcn_permlane16:
     case Intrinsic::amdgcn_permlanex16: {
       // Doing a waterfall loop over these wouldn't make any sense.
       substituteSimpleCopyRegs(OpdMapper, 2);
       substituteSimpleCopyRegs(OpdMapper, 3);
       constrainOpWithReadfirstlane(MI, MRI, 4);
       constrainOpWithReadfirstlane(MI, MRI, 5);
       return;
case Intrinsic::amdgcn_lds_param_load: {
       for (int I = 2; I != M0Idx && MI.getOperand(I).isReg(); ++I)
         OpdsMapping[I] = AMDGPU::getValueMapping(AMDGPU::VGPRRegBankID, 32);
 
       // Must be SGPR, but we must take whatever the original bank is and fix it
       // later.
       OpdsMapping[M0Idx] = AMDGPU::getValueMapping(M0Bank, 32);
       break;
     }
+    case Intrinsic::amdgcn_interp_inreg_p10:
+    case Intrinsic::amdgcn_interp_inreg_p2:
+    case Intrinsic::amdgcn_interp_inreg_p10_f16:
+    case Intrinsic::amdgcn_interp_inreg_p2_f16: {
+      unsigned DstSize = MRI.getType(MI.getOperand(0).getReg()).getSizeInBits();
+      OpdsMapping[0] = AMDGPU::getValueMapping(AMDGPU::VGPRRegBankID, DstSize);
+      OpdsMapping[2] = AMDGPU::getValueMapping(AMDGPU::VGPRRegBankID, 32);
+      OpdsMapping[3] = AMDGPU::getValueMapping(AMDGPU::VGPRRegBankID, 32);
+      OpdsMapping[4] = AMDGPU::getValueMapping(AMDGPU::VGPRRegBankID, 32);
+      break;
+    }
     case Intrinsic::amdgcn_ballot: {
       unsigned DstSize = MRI.getType(MI.getOperand(0).getReg()).getSizeInBits();
       unsigned SrcSize = MRI.getType(MI.getOperand(2).getReg()).getSizeInBits();
       OpdsMapping[0] = AMDGPU::getValueMapping(AMDGPU::SGPRRegBankID, DstSize);
       OpdsMapping[2] = AMDGPU::getValueMapping(AMDGPU::VCCRegBankID, SrcSize);
       break;
     }
     }
case AMDGPU::G_AMDGPU_INTRIN_BVH_INTERSECT_RAY: {
     if (N == 3) {
       // Sequential form: all operands combined into VGPR256/VGPR512
       unsigned Size = MRI.getType(MI.getOperand(2).getReg()).getSizeInBits();
       if (Size > 256)
         Size = 512;
       OpdsMapping[2] = AMDGPU::getValueMapping(AMDGPU::VGPRRegBankID, Size);
     } else {
       // NSA form
-      for (unsigned I = 2; I < N; ++I)
-        OpdsMapping[I] = AMDGPU::getValueMapping(AMDGPU::VGPRRegBankID, 32);
+      for (unsigned I = 2; I < N; ++I) {
+        unsigned Size = MRI.getType(MI.getOperand(I).getReg()).getSizeInBits();
+        OpdsMapping[I] = AMDGPU::getValueMapping(AMDGPU::VGPRRegBankID, Size);
+      }
     }
     break;
   }
   case AMDGPU::G_INTRINSIC_W_SIDE_EFFECTS: {
     auto IntrID = MI.getIntrinsicID();
     switch (IntrID) {
     case Intrinsic::amdgcn_s_getreg:
     case Intrinsic::amdgcn_s_memtime:

llvm/lib/Target/AMDGPU/VINTERPInstructions.td

 def V_INTERP_P10_RTZ_F16_F32_inreg :
   VINTERP_Pseudo <"v_interp_p10_rtz_f16_f32", VOP3_VINTERP_F16<[f32, f32, f32, f32]>>;
 def V_INTERP_P2_RTZ_F16_F32_inreg :
   VINTERP_Pseudo <"v_interp_p2_rtz_f16_f32", VOP3_VINTERP_F16<[f16, f32, f32, f32]>>;
 } // Uses = [M0, EXEC]
 
 } // SubtargetPredicate = isGFX11Plus
 
+class VInterpF32Pat <SDPatternOperator op, Instruction inst> : GCNPat <
+   (f32 (op
+      (VINTERPMods f32:$src0, i32:$src0_modifiers),
+      (VINTERPMods f32:$src1, i32:$src1_modifiers),
+      (VINTERPMods f32:$src2, i32:$src2_modifiers))),
+    (inst $src0_modifiers, $src0,
+          $src1_modifiers, $src1,
+          $src2_modifiers, $src2,
+          0, /* clamp */
+          7) /* wait_exp */
+>;
+
 def VINTERP_OPSEL {
   int LOW = 0;
   int HIGH = 0xa;
 }
+
+class VInterpF16Pat <SDPatternOperator op, Instruction inst,
+                     ValueType dst_type, bit high,
+                     list<ComplexPattern> pat> : GCNPat <
+   (dst_type (op
+      (pat[0] f32:$src0, i32:$src0_modifiers),
+      (pat[1] f32:$src1, i32:$src1_modifiers),
+      (pat[2] f32:$src2, i32:$src2_modifiers),
+      !if(high, (i1 -1), (i1 0)))),
+    (inst $src0_modifiers, $src0,
+          $src1_modifiers, $src1,
+          $src2_modifiers, $src2,
+          0, /* clamp */
+          /* op_sel = 0 */
+          7) /* wait_exp */
+>;
+
+multiclass VInterpF16Pat <SDPatternOperator op, Instruction inst,
+                          ValueType dst_type, list<ComplexPattern> high_pat> {
+  def : VInterpF16Pat<op, inst, dst_type, 0,
+                      [VINTERPMods, VINTERPMods, VINTERPMods]>;
+  def : VInterpF16Pat<op, inst, dst_type, 1, high_pat>;
+}
+
+def : VInterpF32Pat<int_amdgcn_interp_inreg_p10, V_INTERP_P10_F32_inreg>;
+def : VInterpF32Pat<int_amdgcn_interp_inreg_p2, V_INTERP_P2_F32_inreg>;
+defm : VInterpF16Pat<int_amdgcn_interp_inreg_p10_f16,
+                     V_INTERP_P10_F16_F32_inreg, f32,
+                     [VINTERPModsHi, VINTERPMods, VINTERPModsHi]>;
+defm : VInterpF16Pat<int_amdgcn_interp_inreg_p2_f16,
+                     V_INTERP_P2_F16_F32_inreg, f16,
+                     [VINTERPModsHi, VINTERPMods, VINTERPMods]>;
+
 //===----------------------------------------------------------------------===//
 // VINTERP Real Instructions
 //===----------------------------------------------------------------------===//
 
 let AssemblerPredicate = isGFX11Plus, DecoderNamespace = "GFX11" in {
   multiclass VINTERP_Real_gfx11 <bits<7> op> {
     def _gfx11 :
       VINTERP_Real<!cast<VOP3_Pseudo>(NAME), SIEncodingFamily.GFX11>,

llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.interp.inreg.ll

+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
+; RUN: llc -global-isel -march=amdgcn -mcpu=gfx1100 -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
+
+define amdgpu_ps void @v_interp_f32(float inreg %i, float inreg %j, i32 inreg %m0) #0 {
+; GCN-LABEL: v_interp_f32:
+; GCN:       ; %bb.0: ; %main_body
+; GCN-NEXT:    s_mov_b32 m0, s2
+; GCN-NEXT:    v_mov_b32_e32 v2, s0
+; GCN-NEXT:    lds_param_load v0, attr0.y
+; GCN-NEXT:    lds_param_load v1, attr1.x
+; GCN-NEXT:    v_mov_b32_e32 v4, s1
+; GCN-NEXT:    v_interp_p10_f32 v3, v0, v2, v0 wait_exp:7
+; GCN-NEXT:    v_interp_p10_f32 v2, v1, v2, v1 wait_exp:7
+; GCN-NEXT:    v_interp_p2_f32 v0, v0, v4, v3 wait_exp:7
+; GCN-NEXT:    v_interp_p2_f32 v1, v1, v4, v0 wait_exp:7
+; GCN-NEXT:    exp mrt0 v3, v2, v0, v1 done
+; GCN-NEXT:    s_endpgm
+main_body:
+  %p0 = call float @llvm.amdgcn.lds.param.load(i32 1, i32 0, i32 %m0)
+  %p1 = call float @llvm.amdgcn.lds.param.load(i32 0, i32 1, i32 %m0)
+  %p0_0 = call float @llvm.amdgcn.interp.inreg.p10(float %p0, float %i, float %p0)
+  %p1_0 = call float @llvm.amdgcn.interp.inreg.p2(float %p0, float %j, float %p0_0)
+  %p0_1 = call float @llvm.amdgcn.interp.inreg.p10(float %p1, float %i, float %p1)
+  %p1_1 = call float @llvm.amdgcn.interp.inreg.p2(float %p1, float %j, float %p1_0)
+  call void @llvm.amdgcn.exp.f32(i32 0, i32 15, float %p0_0, float %p0_1, float %p1_0, float %p1_1, i1 true, i1 true) #0
+  ret void
+}
+
+define amdgpu_ps void @v_interp_f32_many(float inreg %i, float inreg %j, i32 inreg %m0) #0 {
+; GCN-LABEL: v_interp_f32_many:
+; GCN:       ; %bb.0: ; %main_body
+; GCN-NEXT:    s_mov_b32 m0, s2
+; GCN-NEXT:    v_mov_b32_e32 v4, s0
+; GCN-NEXT:    lds_param_load v0, attr0.x
+; GCN-NEXT:    lds_param_load v1, attr1.x
+; GCN-NEXT:    lds_param_load v2, attr2.x
+; GCN-NEXT:    lds_param_load v3, attr3.x
+; GCN-NEXT:    v_mov_b32_e32 v5, s1
+; GCN-NEXT:    v_interp_p10_f32 v6, v0, v4, v0 wait_exp:7
+; GCN-NEXT:    v_interp_p10_f32 v7, v1, v4, v1 wait_exp:7
+; GCN-NEXT:    v_interp_p10_f32 v8, v2, v4, v2 wait_exp:7
+; GCN-NEXT:    v_interp_p10_f32 v4, v3, v4, v3 wait_exp:7
+; GCN-NEXT:    v_interp_p2_f32 v0, v0, v5, v6 wait_exp:7
+; GCN-NEXT:    v_interp_p2_f32 v1, v1, v5, v7 wait_exp:7
+; GCN-NEXT:    v_interp_p2_f32 v2, v2, v5, v8 wait_exp:7
+; GCN-NEXT:    v_interp_p2_f32 v3, v3, v5, v4 wait_exp:7
+; GCN-NEXT:    exp mrt0 v0, v1, v2, v3 done
+; GCN-NEXT:    s_endpgm
+main_body:
+  %p0 = call float @llvm.amdgcn.lds.param.load(i32 0, i32 0, i32 %m0)
+  %p1 = call float @llvm.amdgcn.lds.param.load(i32 0, i32 1, i32 %m0)
+  %p2 = call float @llvm.amdgcn.lds.param.load(i32 0, i32 2, i32 %m0)
+  %p3 = call float @llvm.amdgcn.lds.param.load(i32 0, i32 3, i32 %m0)
+  %p0_0 = call float @llvm.amdgcn.interp.inreg.p10(float %p0, float %i, float %p0)
+  %p0_1 = call float @llvm.amdgcn.interp.inreg.p2(float %p0, float %j, float %p0_0)
+  %p1_0 = call float @llvm.amdgcn.interp.inreg.p10(float %p1, float %i, float %p1)
+  %p1_1 = call float @llvm.amdgcn.interp.inreg.p2(float %p1, float %j, float %p1_0)
+  %p2_0 = call float @llvm.amdgcn.interp.inreg.p10(float %p2, float %i, float %p2)
+  %p2_1 = call float @llvm.amdgcn.interp.inreg.p2(float %p2, float %j, float %p2_0)
+  %p3_0 = call float @llvm.amdgcn.interp.inreg.p10(float %p3, float %i, float %p3)
+  %p3_1 = call float @llvm.amdgcn.interp.inreg.p2(float %p3, float %j, float %p3_0)
+  call void @llvm.amdgcn.exp.f32(i32 0, i32 15, float %p0_1, float %p1_1, float %p2_1, float %p3_1, i1 true, i1 true) #0
+  ret void
+}
+
+define amdgpu_ps void @v_interp_f32_many_vm(float addrspace(1)* %ptr, i32 inreg %m0) #0 {
+; GCN-LABEL: v_interp_f32_many_vm:
+; GCN:       ; %bb.0: ; %main_body
+; GCN-NEXT:    global_load_b64 v[0:1], v[0:1], off offset:4
+; GCN-NEXT:    s_mov_b32 m0, s0
+; GCN-NEXT:    lds_param_load v2, attr0.x
+; GCN-NEXT:    lds_param_load v3, attr1.x
+; GCN-NEXT:    lds_param_load v4, attr2.x
+; GCN-NEXT:    lds_param_load v5, attr3.x
+; GCN-NEXT:    s_waitcnt vmcnt(0)
+; GCN-NEXT:    v_interp_p10_f32 v6, v2, v0, v2 wait_exp:7
+; GCN-NEXT:    v_interp_p10_f32 v7, v3, v0, v3 wait_exp:7
+; GCN-NEXT:    v_interp_p10_f32 v8, v4, v0, v4 wait_exp:7
+; GCN-NEXT:    v_interp_p10_f32 v0, v5, v0, v5 wait_exp:7
+; GCN-NEXT:    v_interp_p2_f32 v2, v2, v1, v6 wait_exp:7
+; GCN-NEXT:    v_interp_p2_f32 v3, v3, v1, v7 wait_exp:7
+; GCN-NEXT:    v_interp_p2_f32 v4, v4, v1, v8 wait_exp:7
+; GCN-NEXT:    v_interp_p2_f32 v0, v5, v1, v0 wait_exp:7
+; GCN-NEXT:    exp mrt0 v2, v3, v4, v0 done
+; GCN-NEXT:    s_endpgm
+main_body:
+  %i.ptr = getelementptr float, float addrspace(1)* %ptr, i32 1
+  %i = load float, float addrspace(1)* %i.ptr, align 4
+  %j.ptr = getelementptr float, float addrspace(1)* %ptr, i32 2
+  %j = load float, float addrspace(1)* %j.ptr, align 4
+  %p0 = call float @llvm.amdgcn.lds.param.load(i32 0, i32 0, i32 %m0)
+  %p1 = call float @llvm.amdgcn.lds.param.load(i32 0, i32 1, i32 %m0)
+  %p2 = call float @llvm.amdgcn.lds.param.load(i32 0, i32 2, i32 %m0)
+  %p3 = call float @llvm.amdgcn.lds.param.load(i32 0, i32 3, i32 %m0)
+  %p0_0 = call float @llvm.amdgcn.interp.inreg.p10(float %p0, float %i, float %p0)
+  %p0_1 = call float @llvm.amdgcn.interp.inreg.p2(float %p0, float %j, float %p0_0)
+  %p1_0 = call float @llvm.amdgcn.interp.inreg.p10(float %p1, float %i, float %p1)
+  %p1_1 = call float @llvm.amdgcn.interp.inreg.p2(float %p1, float %j, float %p1_0)
+  %p2_0 = call float @llvm.amdgcn.interp.inreg.p10(float %p2, float %i, float %p2)
+  %p2_1 = call float @llvm.amdgcn.interp.inreg.p2(float %p2, float %j, float %p2_0)
+  %p3_0 = call float @llvm.amdgcn.interp.inreg.p10(float %p3, float %i, float %p3)
+  %p3_1 = call float @llvm.amdgcn.interp.inreg.p2(float %p3, float %j, float %p3_0)
+  call void @llvm.amdgcn.exp.f32(i32 0, i32 15, float %p0_1, float %p1_1, float %p2_1, float %p3_1, i1 true, i1 true) #0
+  ret void
+}
+
+define amdgpu_ps half @v_interp_f16(float inreg %i, float inreg %j, i32 inreg %m0) #0 {
+; GCN-LABEL: v_interp_f16:
+; GCN:       ; %bb.0: ; %main_body
+; GCN-NEXT:    s_mov_b32 m0, s2
+; GCN-NEXT:    v_mov_b32_e32 v1, s0
+; GCN-NEXT:    lds_param_load v0, attr0.x
+; GCN-NEXT:    v_mov_b32_e32 v2, s1
+; GCN-NEXT:    v_interp_p10_f16_f32 v3, v0, v1, v0 wait_exp:7
+; GCN-NEXT:    v_interp_p10_f16_f32 v1, v0, v1, v0 op_sel:[1,0,1,0] wait_exp:7
+; GCN-NEXT:    v_interp_p2_f16_f32 v3, v0, v2, v3 wait_exp:7
+; GCN-NEXT:    v_interp_p2_f16_f32 v0, v0, v2, v1 op_sel:[1,0,0,0] wait_exp:7
+; GCN-NEXT:    v_add_f16_e32 v0, v3, v0
+; GCN-NEXT:    ; return to shader part epilog
+main_body:
+  %p0 = call float @llvm.amdgcn.lds.param.load(i32 0, i32 0, i32 %m0)
+  %l_p0 = call float @llvm.amdgcn.interp.inreg.p10.f16(float %p0, float %i, float %p0, i1 0)
+  %l_p1 = call half @llvm.amdgcn.interp.inreg.p2.f16(float %p0, float %j, float %l_p0, i1 0)
+  %h_p0 = call float @llvm.amdgcn.interp.inreg.p10.f16(float %p0, float %i, float %p0, i1 1)
+  %h_p1 = call half @llvm.amdgcn.interp.inreg.p2.f16(float %p0, float %j, float %h_p0, i1 1)
+  %res = fadd half %l_p1, %h_p1
+  ret half %res
+}
+
+declare float @llvm.amdgcn.lds.param.load(i32, i32, i32) #1
+declare float @llvm.amdgcn.interp.inreg.p10(float, float, float) #0
+declare float @llvm.amdgcn.interp.inreg.p2(float, float, float) #0
+declare float @llvm.amdgcn.interp.inreg.p10.f16(float, float, float, i1) #0
+declare half @llvm.amdgcn.interp.inreg.p2.f16(float, float, float, i1) #0
+declare void @llvm.amdgcn.exp.f32(i32, i32, float, float, float, float, i1, i1) #0
+declare void @llvm.amdgcn.exp.f16(i32, i32, float, float, float, float, i1, i1) #0
+
+attributes #0 = { nounwind }
+attributes #1 = { nounwind readnone }

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.interp.inreg.ll

+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
+; RUN: llc -march=amdgcn -mcpu=gfx1100 -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
+
+define amdgpu_ps void @v_interp_f32(float inreg %i, float inreg %j, i32 inreg %m0) #0 {
+; GCN-LABEL: v_interp_f32:
+; GCN:       ; %bb.0: ; %main_body
+; GCN-NEXT:    s_mov_b32 m0, s2
+; GCN-NEXT:    v_mov_b32_e32 v2, s0
+; GCN-NEXT:    lds_param_load v0, attr0.y
+; GCN-NEXT:    lds_param_load v1, attr1.x
+; GCN-NEXT:    v_mov_b32_e32 v4, s1
+; GCN-NEXT:    v_interp_p10_f32 v3, v0, v2, v0 wait_exp:7
+; GCN-NEXT:    v_interp_p10_f32 v2, v1, v2, v1 wait_exp:7
+; GCN-NEXT:    v_interp_p2_f32 v0, v0, v4, v3 wait_exp:7
+; GCN-NEXT:    v_interp_p2_f32 v1, v1, v4, v0 wait_exp:7
+; GCN-NEXT:    exp mrt0 v3, v2, v0, v1 done
+; GCN-NEXT:    s_endpgm
+main_body:
+  %p0 = call float @llvm.amdgcn.lds.param.load(i32 1, i32 0, i32 %m0)
+  %p1 = call float @llvm.amdgcn.lds.param.load(i32 0, i32 1, i32 %m0)
+  %p0_0 = call float @llvm.amdgcn.interp.inreg.p10(float %p0, float %i, float %p0)
+  %p1_0 = call float @llvm.amdgcn.interp.inreg.p2(float %p0, float %j, float %p0_0)
+  %p0_1 = call float @llvm.amdgcn.interp.inreg.p10(float %p1, float %i, float %p1)
+  %p1_1 = call float @llvm.amdgcn.interp.inreg.p2(float %p1, float %j, float %p1_0)
+  call void @llvm.amdgcn.exp.f32(i32 0, i32 15, float %p0_0, float %p0_1, float %p1_0, float %p1_1, i1 true, i1 true) #0
+  ret void
+}
+
+define amdgpu_ps void @v_interp_f32_many(float inreg %i, float inreg %j, i32 inreg %m0) #0 {
+; GCN-LABEL: v_interp_f32_many:
+; GCN:       ; %bb.0: ; %main_body
+; GCN-NEXT:    s_mov_b32 m0, s2
+; GCN-NEXT:    v_mov_b32_e32 v4, s0
+; GCN-NEXT:    lds_param_load v0, attr0.x
+; GCN-NEXT:    lds_param_load v1, attr1.x
+; GCN-NEXT:    lds_param_load v2, attr2.x
+; GCN-NEXT:    lds_param_load v3, attr3.x
+; GCN-NEXT:    v_mov_b32_e32 v5, s1
+; GCN-NEXT:    v_interp_p10_f32 v6, v0, v4, v0 wait_exp:7
+; GCN-NEXT:    v_interp_p10_f32 v7, v1, v4, v1 wait_exp:7
+; GCN-NEXT:    v_interp_p10_f32 v8, v2, v4, v2 wait_exp:7
+; GCN-NEXT:    v_interp_p10_f32 v4, v3, v4, v3 wait_exp:7
+; GCN-NEXT:    v_interp_p2_f32 v0, v0, v5, v6 wait_exp:7
+; GCN-NEXT:    v_interp_p2_f32 v1, v1, v5, v7 wait_exp:7
+; GCN-NEXT:    v_interp_p2_f32 v2, v2, v5, v8 wait_exp:7
+; GCN-NEXT:    v_interp_p2_f32 v3, v3, v5, v4 wait_exp:7
+; GCN-NEXT:    exp mrt0 v0, v1, v2, v3 done
+; GCN-NEXT:    s_endpgm
+main_body:
+  %p0 = call float @llvm.amdgcn.lds.param.load(i32 0, i32 0, i32 %m0)
+  %p1 = call float @llvm.amdgcn.lds.param.load(i32 0, i32 1, i32 %m0)
+  %p2 = call float @llvm.amdgcn.lds.param.load(i32 0, i32 2, i32 %m0)
+  %p3 = call float @llvm.amdgcn.lds.param.load(i32 0, i32 3, i32 %m0)
+  %p0_0 = call float @llvm.amdgcn.interp.inreg.p10(float %p0, float %i, float %p0)
+  %p0_1 = call float @llvm.amdgcn.interp.inreg.p2(float %p0, float %j, float %p0_0)
+  %p1_0 = call float @llvm.amdgcn.interp.inreg.p10(float %p1, float %i, float %p1)
+  %p1_1 = call float @llvm.amdgcn.interp.inreg.p2(float %p1, float %j, float %p1_0)
+  %p2_0 = call float @llvm.amdgcn.interp.inreg.p10(float %p2, float %i, float %p2)
+  %p2_1 = call float @llvm.amdgcn.interp.inreg.p2(float %p2, float %j, float %p2_0)
+  %p3_0 = call float @llvm.amdgcn.interp.inreg.p10(float %p3, float %i, float %p3)
+  %p3_1 = call float @llvm.amdgcn.interp.inreg.p2(float %p3, float %j, float %p3_0)
+  call void @llvm.amdgcn.exp.f32(i32 0, i32 15, float %p0_1, float %p1_1, float %p2_1, float %p3_1, i1 true, i1 true) #0
+  ret void
+}
+
+define amdgpu_ps void @v_interp_f32_many_vm(float addrspace(1)* %ptr, i32 inreg %m0) #0 {
+; GCN-LABEL: v_interp_f32_many_vm:
+; GCN:       ; %bb.0: ; %main_body
+; GCN-NEXT:    global_load_b64 v[0:1], v[0:1], off offset:4
+; GCN-NEXT:    s_mov_b32 m0, s0
+; GCN-NEXT:    lds_param_load v2, attr0.x
+; GCN-NEXT:    lds_param_load v3, attr1.x
+; GCN-NEXT:    lds_param_load v4, attr2.x
+; GCN-NEXT:    lds_param_load v5, attr3.x
+; GCN-NEXT:    s_waitcnt vmcnt(0)
+; GCN-NEXT:    v_interp_p10_f32 v6, v2, v0, v2 wait_exp:7
+; GCN-NEXT:    v_interp_p10_f32 v7, v3, v0, v3 wait_exp:7
+; GCN-NEXT:    v_interp_p10_f32 v8, v4, v0, v4 wait_exp:7
+; GCN-NEXT:    v_interp_p10_f32 v0, v5, v0, v5 wait_exp:7
+; GCN-NEXT:    v_interp_p2_f32 v2, v2, v1, v6 wait_exp:7
+; GCN-NEXT:    v_interp_p2_f32 v3, v3, v1, v7 wait_exp:7
+; GCN-NEXT:    v_interp_p2_f32 v4, v4, v1, v8 wait_exp:7
+; GCN-NEXT:    v_interp_p2_f32 v0, v5, v1, v0 wait_exp:7
+; GCN-NEXT:    exp mrt0 v2, v3, v4, v0 done
+; GCN-NEXT:    s_endpgm
+main_body:
+  %i.ptr = getelementptr float, float addrspace(1)* %ptr, i32 1
+  %i = load float, float addrspace(1)* %i.ptr, align 4
+  %j.ptr = getelementptr float, float addrspace(1)* %ptr, i32 2
+  %j = load float, float addrspace(1)* %j.ptr, align 4
+  %p0 = call float @llvm.amdgcn.lds.param.load(i32 0, i32 0, i32 %m0)
+  %p1 = call float @llvm.amdgcn.lds.param.load(i32 0, i32 1, i32 %m0)
+  %p2 = call float @llvm.amdgcn.lds.param.load(i32 0, i32 2, i32 %m0)
+  %p3 = call float @llvm.amdgcn.lds.param.load(i32 0, i32 3, i32 %m0)
+  %p0_0 = call float @llvm.amdgcn.interp.inreg.p10(float %p0, float %i, float %p0)
+  %p0_1 = call float @llvm.amdgcn.interp.inreg.p2(float %p0, float %j, float %p0_0)
+  %p1_0 = call float @llvm.amdgcn.interp.inreg.p10(float %p1, float %i, float %p1)
+  %p1_1 = call float @llvm.amdgcn.interp.inreg.p2(float %p1, float %j, float %p1_0)
+  %p2_0 = call float @llvm.amdgcn.interp.inreg.p10(float %p2, float %i, float %p2)
+  %p2_1 = call float @llvm.amdgcn.interp.inreg.p2(float %p2, float %j, float %p2_0)
+  %p3_0 = call float @llvm.amdgcn.interp.inreg.p10(float %p3, float %i, float %p3)
+  %p3_1 = call float @llvm.amdgcn.interp.inreg.p2(float %p3, float %j, float %p3_0)
+  call void @llvm.amdgcn.exp.f32(i32 0, i32 15, float %p0_1, float %p1_1, float %p2_1, float %p3_1, i1 true, i1 true) #0
+  ret void
+}
+
+define amdgpu_ps half @v_interp_f16(float inreg %i, float inreg %j, i32 inreg %m0) #0 {
+; GCN-LABEL: v_interp_f16:
+; GCN:       ; %bb.0: ; %main_body
+; GCN-NEXT:    s_mov_b32 m0, s2
+; GCN-NEXT:    v_mov_b32_e32 v1, s0
+; GCN-NEXT:    lds_param_load v0, attr0.x
+; GCN-NEXT:    v_mov_b32_e32 v2, s1
+; GCN-NEXT:    v_interp_p10_f16_f32 v3, v0, v1, v0 wait_exp:7
+; GCN-NEXT:    v_interp_p10_f16_f32 v1, v0, v1, v0 op_sel:[1,0,1,0] wait_exp:7
+; GCN-NEXT:    v_interp_p2_f16_f32 v3, v0, v2, v3 wait_exp:7
+; GCN-NEXT:    v_interp_p2_f16_f32 v0, v0, v2, v1 op_sel:[1,0,0,0] wait_exp:7
+; GCN-NEXT:    v_add_f16_e32 v0, v3, v0
+; GCN-NEXT:    ; return to shader part epilog
+main_body:
+  %p0 = call float @llvm.amdgcn.lds.param.load(i32 0, i32 0, i32 %m0)
+  %l_p0 = call float @llvm.amdgcn.interp.inreg.p10.f16(float %p0, float %i, float %p0, i1 0)
+  %l_p1 = call half @llvm.amdgcn.interp.inreg.p2.f16(float %p0, float %j, float %l_p0, i1 0)
+  %h_p0 = call float @llvm.amdgcn.interp.inreg.p10.f16(float %p0, float %i, float %p0, i1 1)
+  %h_p1 = call half @llvm.amdgcn.interp.inreg.p2.f16(float %p0, float %j, float %h_p0, i1 1)
+  %res = fadd half %l_p1, %h_p1
+  ret half %res
+}
+
+declare float @llvm.amdgcn.lds.param.load(i32, i32, i32) #1
+declare float @llvm.amdgcn.interp.inreg.p10(float, float, float) #0
+declare float @llvm.amdgcn.interp.inreg.p2(float, float, float) #0
+declare float @llvm.amdgcn.interp.inreg.p10.f16(float, float, float, i1) #0
+declare half @llvm.amdgcn.interp.inreg.p2.f16(float, float, float, i1) #0
+declare void @llvm.amdgcn.exp.f32(i32, i32, float, float, float, float, i1, i1) #0
+declare void @llvm.amdgcn.exp.f16(i32, i32, float, float, float, float, i1, i1) #0
+
+attributes #0 = { nounwind }
+attributes #1 = { nounwind readnone }