[AMDGPU] gfx11 ldsdir intrinsics and ISel

[foad]

Can we also add something like test/CodeGen/AMDGPU/llvm.amdgcn.lds.direct.load.ll and test/CodeGen/AMDGPU/llvm.amdgcn.lds.param.load.ll to test codegen?

[Joe_Nash]

added ISA codegen test.

[Joe_Nash]

In D127664#3578415, @foad wrote:

Can we also add something like test/CodeGen/AMDGPU/llvm.amdgcn.lds.direct.load.ll and test/CodeGen/AMDGPU/llvm.amdgcn.lds.param.load.ll to test codegen?

I have added something. Just to note, changes to SIInsertWaitcnt are coming which will affect these tests.

[arsenm]

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

1464

Why is the return type hardcoded to float instead of mangled for a load?

1465

Also would expect this to be an addrspace 3 pointer

[critson]

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

1465

This is setting M0, the same as parameter loads.
Construction of M0 values is handled by the front-end because it is not a pure address pointer but rather a combination of address offset and flags describing the data type.
I guess we could rework this to form M0 in the backend based on an address 3 pointer and a return type.

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.lds.param.load.ll

3 ↗

(On Diff #436442)

There should probably be a move to m0 here?

[nhaehnle]

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

1465

I think it makes sense to have an intrinsic that closely corresponds to the instruction itself. Maybe add a comment explaining this fact about M0?

I agree with Matt that the return value should be mangled.

[Joe_Nash]

changed return type of lds_direct_load and commented on input argument type

[Joe_Nash]

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

1465

I have commented on the input argument, and changed the return type. I am not familiar with the type mangling here, does this look correct?

[Joe_Nash]

added s_mov m0 to param.load test

[Joe_Nash]

removed builtin from intrinsic definition

[Joe_Nash]

I think all outstanding issues are addressed; please take another look.

[arsenm]

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.lds.direct.load.ll

2 ↗

(On Diff #437327)

Missing globalisel codegen tests

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.lds.param.load.ll

2 ↗

(On Diff #437327)

Missing globalisel codegen tests

[Joe_Nash]

added globalisel runlines to codegen test

[rampitec]

llvm/lib/Target/AMDGPU/LDSDIRInstructions.td

96

How does that work? It seems to ignore M0 argument.

[Joe_Nash]

llvm/lib/Target/AMDGPU/LDSDIRInstructions.td

96

Because LDS_DIRECT_LOAD has an implicit use of M0. In llvm.amdgcn.lds.direct.load.ll it seems to work as it should given that.

[rampitec]

llvm/lib/Target/AMDGPU/LDSDIRInstructions.td

96

Right, it uses M0, but where is a link from the call argument and actual store of that value into M0?

[Joe_Nash]

llvm/lib/Target/AMDGPU/LDSDIRInstructions.td

96

I don't know, but maybe this helps to explain? Dump from AMDGPUGenGlobalISel.inc

// Label 2092: @108525
GIM_Try, /*On fail goto*//*Label 2093*/ 108577, // Rule ID 3516 //
  GIM_CheckIntrinsicID, /*MI*/0, /*Op*/1, Intrinsic::amdgcn_lds_direct_load,
  GIM_CheckType, /*MI*/0, /*Op*/0, /*Type*/GILLT_s32,
  GIM_CheckType, /*MI*/0, /*Op*/2, /*Type*/GILLT_s32,
  GIM_CheckRegBankForClass, /*MI*/0, /*Op*/0, /*RC*/AMDGPU::VGPR_32RegClassID,
  GIM_CheckRegBankForClass, /*MI*/0, /*Op*/2, /*RC*/AMDGPU::M0_CLASSRegClassID,
  // (intrinsic_w_chain:{ *:[f32] } 1864:{ *:[iPTR] }, M0:{ *:[i32] })  =>  (LDS_DIRECT_LOAD:{ *:[f32] } 0:{ *:[i8] })
  GIR_BuildMI, /*InsnID*/1, /*Opcode*/TargetOpcode::COPY,
  GIR_AddRegister, /*InsnID*/1, AMDGPU::M0, /*AddRegisterRegFlags*/RegState::Define,
  GIR_Copy, /*NewInsnID*/1, /*OldInsnID*/0, /*OpIdx*/2, // M0
  GIR_BuildMI, /*InsnID*/0, /*Opcode*/AMDGPU::LDS_DIRECT_LOAD,
  GIR_Copy, /*NewInsnID*/0, /*OldInsnID*/0, /*OpIdx*/0, // vdst
  GIR_AddImm, /*InsnID*/0, /*Imm*/0,
  GIR_MergeMemOperands, /*InsnID*/0, /*MergeInsnID's*/0, GIU_MergeMemOperands_EndOfList,
  GIR_EraseFromParent, /*InsnID*/0,
  GIR_ConstrainSelectedInstOperands, /*InsnID*/0,
  // GIR_Coverage, 3516,
  GIR_Done,
// Label 2093: @108577
GIM_Reject,

[rampitec]

LGTM

llvm/lib/Target/AMDGPU/LDSDIRInstructions.td

96

OK, I must admit I still do not understand how does it work, but it obviously does.

[arsenm]

llvm/lib/Target/AMDGPU/LDSDIRInstructions.td

96

The physical register takes the place in the pattern of what normally would be a named virtual input, like VReg_32:$src1

[foad]

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

1461

The __int_* prefix doesn't make much sense. I would suggest either using the tablegen name (int_amdgcn_lds_direct_load) or preferably the LLVM IR name (llvm.amdgcn.lds.direct.load).

[Joe_Nash]

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

1461

Done, see
75378d432fda5408b7210fd3627db884561db650

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

 // high selects whether high or low 16-bits are loaded from LDS
 def int_amdgcn_interp_p2_f16 :
   GCCBuiltin<"__builtin_amdgcn_interp_p2_f16">,
   Intrinsic<[llvm_half_ty],
             [llvm_float_ty, llvm_float_ty, llvm_i32_ty, llvm_i32_ty, llvm_i1_ty, llvm_i32_ty],
             [IntrNoMem, IntrSpeculatable, IntrWillReturn,
              ImmArg<ArgIndex<2>>, ImmArg<ArgIndex<3>>, ImmArg<ArgIndex<4>>]>;
 
+// __builtin_amdgcn_lds_direct_load <m0>
+def int_amdgcn_lds_direct_load :
+  GCCBuiltin<"__builtin_amdgcn_lds_direct_load">,
+  Intrinsic<[llvm_float_ty],
+            [llvm_i32_ty],
+            [IntrReadMem, IntrSpeculatable, IntrWillReturn]>;
+
+// __builtin_amdgcn_lds_param_load <attr_chan>, <attr>, <m0>
+// Like interp intrinsics, this reads from lds, but the memory values are constant,
+// 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>>]>;
+
 // 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/AMDGPURegisterBankInfo.cpp

case Intrinsic::amdgcn_writelane: {
       constrainOpWithReadfirstlane(MI, MRI, 2); // Source value
       constrainOpWithReadfirstlane(MI, MRI, 3); // Index
       return;
     }
     case Intrinsic::amdgcn_interp_p1:
     case Intrinsic::amdgcn_interp_p2:
     case Intrinsic::amdgcn_interp_mov:
     case Intrinsic::amdgcn_interp_p1_f16:
-    case Intrinsic::amdgcn_interp_p2_f16: {
+    case Intrinsic::amdgcn_interp_p2_f16:
+    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_permlane16:
case Intrinsic::amdgcn_struct_buffer_load_lds: {
       constrainOpWithReadfirstlane(MI, MRI, 6); // soffset
       return;
     }
     case Intrinsic::amdgcn_global_load_lds: {
       applyDefaultMapping(OpdMapper);
       constrainOpWithReadfirstlane(MI, MRI, 2);
       return;
     }
+    case Intrinsic::amdgcn_lds_direct_load: {
+      applyDefaultMapping(OpdMapper);
+      // Readlane for m0 value, which is always the last operand.
+      constrainOpWithReadfirstlane(MI, MRI, MI.getNumOperands() - 1); // Index
+      return;
+    }
     default: {
       if (const AMDGPU::RsrcIntrinsic *RSrcIntrin =
               AMDGPU::lookupRsrcIntrinsic(IntrID)) {
         // Non-images can have complications from operands that allow both SGPR
         // and VGPR. For now it's too complicated to figure out the final opcode
         // to derive the register bank from the MCInstrDesc.
         if (RSrcIntrin->IsImage) {
           applyMappingImage(MI, OpdMapper, MRI, RSrcIntrin->RsrcArg);
case Intrinsic::amdgcn_smfmac_i32_32x32x32_i8: {
       OpdsMapping[4] = getAGPROpMapping(MI.getOperand(4).getReg(), MRI, *TRI);
       OpdsMapping[5] = getVGPROpMapping(MI.getOperand(5).getReg(), MRI, *TRI);
       break;
     }
     case Intrinsic::amdgcn_interp_p1:
     case Intrinsic::amdgcn_interp_p2:
     case Intrinsic::amdgcn_interp_mov:
     case Intrinsic::amdgcn_interp_p1_f16:
-    case Intrinsic::amdgcn_interp_p2_f16: {
+    case Intrinsic::amdgcn_interp_p2_f16:
+    case Intrinsic::amdgcn_lds_param_load: {
       const int M0Idx = MI.getNumOperands() - 1;
       Register M0Reg = MI.getOperand(M0Idx).getReg();
       unsigned M0Bank = getRegBankID(M0Reg, MRI, AMDGPU::SGPRRegBankID);
       unsigned DstSize = MRI.getType(MI.getOperand(0).getReg()).getSizeInBits();
 
       OpdsMapping[0] = AMDGPU::getValueMapping(AMDGPU::VGPRRegBankID, DstSize);
       for (int I = 2; I != M0Idx && MI.getOperand(I).isReg(); ++I)
         OpdsMapping[I] = AMDGPU::getValueMapping(AMDGPU::VGPRRegBankID, 32);
case Intrinsic::amdgcn_ds_gws_sema_release_all: {
       OpdsMapping[1] = AMDGPU::getValueMapping(Bank, 32);
       break;
     }
     case Intrinsic::amdgcn_global_load_lds: {
       OpdsMapping[1] = getVGPROpMapping(MI.getOperand(1).getReg(), MRI, *TRI);
       OpdsMapping[2] = getSGPROpMapping(MI.getOperand(2).getReg(), MRI, *TRI);
       break;
     }
+    case Intrinsic::amdgcn_lds_direct_load: {
+      const int M0Idx = MI.getNumOperands() - 1;
+      Register M0Reg = MI.getOperand(M0Idx).getReg();
+      unsigned M0Bank = getRegBankID(M0Reg, MRI, AMDGPU::SGPRRegBankID);
+      unsigned DstSize = MRI.getType(MI.getOperand(0).getReg()).getSizeInBits();
+
+      OpdsMapping[0] = AMDGPU::getValueMapping(AMDGPU::VGPRRegBankID, DstSize);
+      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;
+    }
     default:
       return getInvalidInstructionMapping();
     }
     break;
   }
   case AMDGPU::G_SELECT: {
     unsigned Size = MRI.getType(MI.getOperand(0).getReg()).getSizeInBits();
     unsigned Op2Bank = getRegBankID(MI.getOperand(2).getReg(), MRI,

llvm/lib/Target/AMDGPU/AMDGPUSearchableTables.td

 def : SourceOfDivergence<int_amdgcn_workitem_id_x>;
 def : SourceOfDivergence<int_amdgcn_workitem_id_y>;
 def : SourceOfDivergence<int_amdgcn_workitem_id_z>;
 def : SourceOfDivergence<int_amdgcn_interp_mov>;
 def : SourceOfDivergence<int_amdgcn_interp_p1>;
 def : SourceOfDivergence<int_amdgcn_interp_p2>;
 def : SourceOfDivergence<int_amdgcn_interp_p1_f16>;
 def : SourceOfDivergence<int_amdgcn_interp_p2_f16>;
+def : SourceOfDivergence<int_amdgcn_lds_direct_load>;
+def : SourceOfDivergence<int_amdgcn_lds_param_load>;
 def : SourceOfDivergence<int_amdgcn_mbcnt_hi>;
 def : SourceOfDivergence<int_amdgcn_mbcnt_lo>;
 def : SourceOfDivergence<int_r600_read_tidig_x>;
 def : SourceOfDivergence<int_r600_read_tidig_y>;
 def : SourceOfDivergence<int_r600_read_tidig_z>;
 def : SourceOfDivergence<int_amdgcn_atomic_inc>;
 def : SourceOfDivergence<int_amdgcn_atomic_dec>;
 def : SourceOfDivergence<int_amdgcn_global_atomic_csub>;

llvm/lib/Target/AMDGPU/LDSDIRInstructions.td

 
 //===----------------------------------------------------------------------===//
 // LDS Direct Instructions
 //===----------------------------------------------------------------------===//
 
 def LDS_DIRECT_LOAD : LDSDIR_Pseudo<"lds_direct_load", 1>;
 def LDS_PARAM_LOAD : LDSDIR_Pseudo<"lds_param_load", 0>;
 
+def : GCNPat <
+  (f32 (int_amdgcn_lds_direct_load M0)),
+  (LDS_DIRECT_LOAD 0)
+>;
+
+def : GCNPat <
+  (f32 (int_amdgcn_lds_param_load timm:$attrchan, timm:$attr, M0)),
+  (LDS_PARAM_LOAD timm:$attr, timm:$attrchan, 0)
+>;
+
 //===----------------------------------------------------------------------===//
 // GFX11+
 //===----------------------------------------------------------------------===//
 
 multiclass LDSDIR_Real_gfx11<bits<2> op, LDSDIR_Pseudo lds = !cast<LDSDIR_Pseudo>(NAME)> {
   def _gfx11 : LDSDIR_Real<op, lds, SIEncodingFamily.GFX11> {
     let AssemblerPredicate = isGFX11Plus;
     let DecoderNamespace = "GFX11";
   }
 }
 
 defm LDS_PARAM_LOAD : LDSDIR_Real_gfx11<0x0>;
 defm LDS_DIRECT_LOAD : LDSDIR_Real_gfx11<0x1>;

llvm/test/CodeGen/AMDGPU/GlobalISel/regbankselect-amdgcn.lds.direct.load.mir

+# NOTE: Assertions have been autogenerated by utils/update_mir_test_checks.py
+# RUN: llc -march=amdgcn -mcpu=gfx1100 -run-pass=regbankselect -regbankselect-fast -verify-machineinstrs -o - %s | FileCheck %s
+# RUN: llc -march=amdgcn -mcpu=gfx1100 -run-pass=regbankselect -regbankselect-greedy -verify-machineinstrs -o - %s | FileCheck %s
+
+---
+name: lds_direct_load_s
+legalized: true
+tracksRegLiveness: true
+
+body: |
+  bb.0:
+    liveins: $sgpr0
+    ; CHECK-LABEL: name: lds_direct_load_s
+    ; CHECK: liveins: $sgpr0
+    ; CHECK: [[COPY:%[0-9]+]]:sgpr(s32) = COPY $sgpr0
+    ; CHECK: [[INT:%[0-9]+]]:vgpr(s32) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.amdgcn.lds.direct.load), [[COPY]](s32)
+    %0:_(s32) = COPY $sgpr0
+    %1:_(s32) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.amdgcn.lds.direct.load), %0
+...
+
+---
+name: lds_direct_load_v
+legalized: true
+tracksRegLiveness: true
+
+body: |
+  bb.0:
+    liveins: $vgpr0
+    ; CHECK-LABEL: name: lds_direct_load_v
+    ; CHECK: liveins: $vgpr0
+    ; CHECK: [[COPY:%[0-9]+]]:vgpr_32(s32) = COPY $vgpr0
+    ; CHECK: [[V_READFIRSTLANE_B32_:%[0-9]+]]:sreg_32(s32) = V_READFIRSTLANE_B32 [[COPY]](s32), implicit $exec
+    ; CHECK: [[INT:%[0-9]+]]:vgpr(s32) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.amdgcn.lds.direct.load), [[V_READFIRSTLANE_B32_]](s32)
+    %0:_(s32) = COPY $vgpr0
+    %1:_(s32) = G_INTRINSIC_W_SIDE_EFFECTS intrinsic(@llvm.amdgcn.lds.direct.load), %0
+...

llvm/test/CodeGen/AMDGPU/GlobalISel/regbankselect-amdgcn.lds.param.load.mir

+# NOTE: Assertions have been autogenerated by utils/update_mir_test_checks.py
+# RUN: llc -march=amdgcn -mcpu=gfx1100 -run-pass=regbankselect -regbankselect-fast -verify-machineinstrs -o - %s | FileCheck %s
+# RUN: llc -march=amdgcn -mcpu=gfx1100 -run-pass=regbankselect -regbankselect-greedy -verify-machineinstrs -o - %s | FileCheck %s
+
+---
+name: lds_param_load_s
+legalized: true
+tracksRegLiveness: true
+
+body: |
+  bb.0:
+    liveins: $sgpr0
+    ; CHECK-LABEL: name: lds_param_load_s
+    ; CHECK: liveins: $sgpr0
+    ; CHECK: [[COPY:%[0-9]+]]:sgpr(s32) = COPY $sgpr0
+    ; CHECK: [[INT:%[0-9]+]]:vgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.lds.param.load), 1, 1, [[COPY]](s32)
+    %0:_(s32) = COPY $sgpr0
+    %1:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.lds.param.load), 1, 1, %0
+...
+
+---
+name: lds_param_load_v
+legalized: true
+tracksRegLiveness: true
+
+body: |
+  bb.0:
+    liveins: $vgpr0
+    ; CHECK-LABEL: name: lds_param_load_v
+    ; CHECK: liveins: $vgpr0
+    ; CHECK: [[COPY:%[0-9]+]]:vgpr_32(s32) = COPY $vgpr0
+    ; CHECK: [[V_READFIRSTLANE_B32_:%[0-9]+]]:sreg_32(s32) = V_READFIRSTLANE_B32 [[COPY]](s32), implicit $exec
+    ; CHECK: [[INT:%[0-9]+]]:vgpr(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.lds.param.load), 1, 1, [[V_READFIRSTLANE_B32_]](s32)
+    %0:_(s32) = COPY $vgpr0
+    %1:_(s32) = G_INTRINSIC intrinsic(@llvm.amdgcn.lds.param.load), 1, 1, %0
+...