AMDGPU: Add s_dcache_* instructions

[• tstellarAMD]

LGTM.

[arsenm]

r248533

include/llvm/IR/IntrinsicsAMDGPU.td

 def int_amdgcn_buffer_wbinvl1_vol :
   GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1_vol">,
   Intrinsic<[], [], []>;
 
 def int_amdgcn_buffer_wbinvl1 :
   GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1">,
   Intrinsic<[], [], []>;
 
+def int_amdgcn_s_dcache_inv :
+  GCCBuiltin<"__builtin_amdgcn_s_dcache_inv">,
+  Intrinsic<[], [], []>;
+
+// CI+
+def int_amdgcn_s_dcache_inv_vol :
+  GCCBuiltin<"__builtin_amdgcn_s_dcache_inv_vol">,
+  Intrinsic<[], [], []>;
+
+// VI
+def int_amdgcn_s_dcache_wb :
+  GCCBuiltin<"__builtin_amdgcn_s_dcache_wb">,
+  Intrinsic<[], [], []>;
+
+// VI
+def int_amdgcn_s_dcache_wb_vol :
+  GCCBuiltin<"__builtin_amdgcn_s_dcache_wb_vol">,
+  Intrinsic<[], [], []>;
+
 }

lib/Target/AMDGPU/CIInstructions.td

 defm V_LOG_LEGACY_F32 : VOP1Inst <vop1<0x45, 0x4c>, "v_log_legacy_f32",
   VOP_F32_F32
 >;
 defm V_EXP_LEGACY_F32 : VOP1Inst <vop1<0x46, 0x4b>, "v_exp_legacy_f32",
   VOP_F32_F32
 >;
 
 //===----------------------------------------------------------------------===//
+// SMRD Instructions
+//===----------------------------------------------------------------------===//
+
+defm S_DCACHE_INV_VOL : SMRD_Inval <smrd<0x1d, 0x22>,
+  "s_dcache_inv_vol", int_amdgcn_s_dcache_inv_vol>;
+
+//===----------------------------------------------------------------------===//
 // MUBUF Instructions
 //===----------------------------------------------------------------------===//
 
 defm BUFFER_WBINVL1_VOL : MUBUF_Invalidate <mubuf<0x70, 0x3f>,
   "buffer_wbinvl1_vol", int_amdgcn_buffer_wbinvl1_vol
 >;
 
 //===----------------------------------------------------------------------===//

lib/Target/AMDGPU/SIInsertWaits.cpp

 
 FunctionPass *llvm::createSIInsertWaits(TargetMachine &tm) {
   return new SIInsertWaits(tm);
 }
 
 Counters SIInsertWaits::getHwCounts(MachineInstr &MI) {
 
   uint64_t TSFlags = TII->get(MI.getOpcode()).TSFlags;
-  Counters Result;
+  Counters Result = { { 0, 0, 0 } };
 
   Result.Named.VM = !!(TSFlags & SIInstrFlags::VM_CNT);
 
   // Only consider stores or EXP for EXP_CNT
   Result.Named.EXP = !!(TSFlags & SIInstrFlags::EXP_CNT &&
       (MI.getOpcode() == AMDGPU::EXP || MI.getDesc().mayStore()));
 
   // LGKM may uses larger values
   if (TSFlags & SIInstrFlags::LGKM_CNT) {
 
     if (TII->isSMRD(MI.getOpcode())) {
 
+      if (MI.getNumOperands() != 0) {
         MachineOperand &Op = MI.getOperand(0);
         assert(Op.isReg() && "First LGKM operand must be a register!");
 
         unsigned Reg = Op.getReg();
+
+        // XXX - What if this is a write into a super register?
         unsigned Size = TRI->getMinimalPhysRegClass(Reg)->getSize();
         Result.Named.LGKM = Size > 4 ? 2 : 1;
-
+      } else {
+        // s_dcache_inv etc. do not have a a destination register. Assume we
+        // want a wait on these.
+        // XXX - What is the right value?
+        Result.Named.LGKM = 1;
+      }
     } else {
       // DS
       Result.Named.LGKM = 1;
     }
 
   } else {
     Result.Named.LGKM = 0;
   }

lib/Target/AMDGPU/SIInstrInfo.td

 }
 
 class sopk <bits<5> si, bits<5> vi = si> {
   field bits<5> SI = si;
   field bits<5> VI = vi;
 }
 
 // Specify an SMRD opcode for SI and SMEM opcode for VI
-class smrd<bits<5> si, bits<5> vi = si> {
-  field bits<5> SI = si;
-  field bits<8> VI = { 0, 0, 0, vi };
+
+// FIXME: This should really be bits<5> si, Tablegen crashes if
+// parameter default value is other parameter with different bit size
+class smrd<bits<8> si, bits<8> vi = si> {
+  field bits<5> SI = si{4-0};
+  field bits<8> VI = vi;
 }
 
 // Execpt for the NONE field, this must be kept in sync with the SISubtarget enum
 // in AMDGPUInstrInfo.cpp
 def SISubtarget {
   int NONE = -1;
   int SI = 0;
   int VI = 1;
class SMRD_Real_si <bits<5> op, string opName, bit imm, dag outs, dag ins,
                     string asm> :
   SMRD <outs, ins, asm, []>,
   SMRDe <op, imm>,
   SIMCInstr<opName, SISubtarget.SI> {
   let AssemblerPredicates = [isSICI];
 }
 
 class SMRD_Real_vi <bits<8> op, string opName, bit imm, dag outs, dag ins,
-                    string asm> :
-  SMRD <outs, ins, asm, []>,
+                    string asm, list<dag> pattern = []> :
+  SMRD <outs, ins, asm, pattern>,
   SMEMe_vi <op, imm>,
   SIMCInstr<opName, SISubtarget.VI> {
   let AssemblerPredicates = [isVI];
 }
 
 multiclass SMRD_m <smrd op, string opName, bit imm, dag outs, dag ins,
                    string asm, list<dag> pattern> {
 
   def "" : SMRD_Pseudo <opName, outs, ins, pattern>;
 
   def _si : SMRD_Real_si <op.SI, opName, imm, outs, ins, asm>;
 
   // glc is only applicable to scalar stores, which are not yet
   // implemented.
   let glc = 0 in {
     def _vi : SMRD_Real_vi <op.VI, opName, imm, outs, ins, asm>;
   }
 }
 
+multiclass SMRD_Inval <smrd op, string opName,
+                       SDPatternOperator node> {
+  let hasSideEffects = 1, mayStore = 1 in {
+    def "" : SMRD_Pseudo <opName, (outs), (ins), [(node)]>;
+
+    let sbase = 0, offset = 0 in {
+      let sdst = 0 in {
+        def _si : SMRD_Real_si <op.SI, opName, 0, (outs), (ins), opName>;
+      }
+
+      let glc = 0, sdata = 0 in {
+        def _vi : SMRD_Real_vi <op.VI, opName, 0, (outs), (ins), opName>;
+      }
+    }
+  }
+}
+
+class SMEM_Inval <bits<8> op, string opName, SDPatternOperator node> :
+  SMRD_Real_vi<op, opName, 0, (outs), (ins), opName, [(node)]> {
+  let hasSideEffects = 1;
+  let mayStore = 1;
+  let sbase = 0;
+  let sdata = 0;
+  let glc = 0;
+  let offset = 0;
+}
+
 multiclass SMRD_Helper <smrd op, string opName, RegisterClass baseClass,
                         RegisterClass dstClass> {
   defm _IMM : SMRD_m <
     op, opName#"_IMM", 1, (outs dstClass:$dst),
     (ins baseClass:$sbase, smrd_offset:$offset),
     opName#" $dst, $sbase, $offset", []
   >;
 

lib/Target/AMDGPU/SIInstructions.td

 
 defm S_BUFFER_LOAD_DWORDX16 : SMRD_Helper <
   smrd<0x0c>, "s_buffer_load_dwordx16", SReg_128, SReg_512
 >;
 
 } // mayLoad = 1
 
 //def S_MEMTIME : SMRD_ <0x0000001e, "s_memtime", []>;
-//def S_DCACHE_INV : SMRD_ <0x0000001f, "s_dcache_inv", []>;
+
+defm S_DCACHE_INV : SMRD_Inval <smrd<0x1f, 0x20>, "s_dcache_inv",
+  int_amdgcn_s_dcache_inv>;
 
 //===----------------------------------------------------------------------===//
 // SOP1 Instructions
 //===----------------------------------------------------------------------===//
 
 let isMoveImm = 1 in {
   let isReMaterializable = 1, isAsCheapAsAMove = 1 in {
     defm S_MOV_B32 : SOP1_32 <sop1<0x03, 0x00>, "s_mov_b32", []>;

 } // End isCommutable = 1
 
 // Remaining instructions:
 // FLAT_*
 // S_CBRANCH_CDBGUSER
 // S_CBRANCH_CDBGSYS
 // S_CBRANCH_CDBGSYS_OR_USER
 // S_CBRANCH_CDBGSYS_AND_USER
-// S_DCACHE_INV_VOL
 // DS_NOP
 // DS_GWS_SEMA_RELEASE_ALL
 // DS_WRAP_RTN_B32
 // DS_CNDXCHG32_RTN_B64
 // DS_WRITE_B96
 // DS_WRITE_B128
 // DS_CONDXCHG32_RTN_B128
 // DS_READ_B96

lib/Target/AMDGPU/VIInstructions.td

 }
 
 def : SI2_VI3Alias <"v_ldexp_f32", V_LDEXP_F32_e64_vi>;
 def : SI2_VI3Alias <"v_cvt_pkaccum_u8_f32", V_CVT_PKACCUM_U8_F32_e64_vi>;
 def : SI2_VI3Alias <"v_cvt_pknorm_i16_f32", V_CVT_PKNORM_I16_F32_e64_vi>;
 def : SI2_VI3Alias <"v_cvt_pknorm_u16_f32", V_CVT_PKNORM_U16_F32_e64_vi>;
 def : SI2_VI3Alias <"v_cvt_pkrtz_f16_f32", V_CVT_PKRTZ_F16_F32_e64_vi>;
 
+//===----------------------------------------------------------------------===//
+// SMEM Instructions
+//===----------------------------------------------------------------------===//
+
+def S_DCACHE_WB : SMEM_Inval <0x21,
+  "s_dcache_wb", int_amdgcn_s_dcache_wb>;
+
+def S_DCACHE_WB_VOL : SMEM_Inval <0x23,
+  "s_dcache_wb_vol", int_amdgcn_s_dcache_wb_vol>;
+
 } // End SIAssemblerPredicate = DisableInst, SubtargetPredicate = isVI
 
 //===----------------------------------------------------------------------===//
 // SMEM Patterns
 //===----------------------------------------------------------------------===//
 
 let Predicates = [isVI] in {
 

test/CodeGen/AMDGPU/llvm.amdgcn.s.dcache.inv.ll

+; RUN: llc -march=amdgcn -mcpu=tahiti -show-mc-encoding < %s | FileCheck -check-prefix=GCN -check-prefix=SI %s
+; RUN: llc -march=amdgcn -mcpu=fiji -show-mc-encoding < %s | FileCheck -check-prefix=GCN -check-prefix=VI %s
+
+declare void @llvm.amdgcn.s.dcache.inv() #0
+
+; GCN-LABEL: {{^}}test_s_dcache_inv:
+; GCN-NEXT: ; BB#0:
+; SI-NEXT: s_dcache_inv ; encoding: [0x00,0x00,0xc0,0xc7]
+; VI-NEXT: s_dcache_inv ; encoding: [0x00,0x00,0x80,0xc0,0x00,0x00,0x00,0x00]
+; GCN-NEXT: s_endpgm
+define void @test_s_dcache_inv() #0 {
+  call void @llvm.amdgcn.s.dcache.inv()
+  ret void
+}
+
+; GCN-LABEL: {{^}}test_s_dcache_inv_insert_wait:
+; GCN-NEXT: ; BB#0:
+; GCN-NEXT: s_dcache_inv
+; GCN-NEXT: s_waitcnt lgkmcnt(0) ; encoding
+define void @test_s_dcache_inv_insert_wait() #0 {
+  call void @llvm.amdgcn.s.dcache.inv()
+  br label %end
+
+end:
+  store volatile i32 3, i32 addrspace(1)* undef
+  ret void
+}
+
+attributes #0 = { nounwind }

test/CodeGen/AMDGPU/llvm.amdgcn.s.dcache.inv.vol.ll

+; RUN: llc -march=amdgcn -mcpu=bonaire -show-mc-encoding < %s | FileCheck -check-prefix=GCN -check-prefix=CI %s
+; RUN: llc -march=amdgcn -mcpu=tonga -show-mc-encoding < %s | FileCheck -check-prefix=GCN -check-prefix=VI %s
+
+declare void @llvm.amdgcn.s.dcache.inv.vol() #0
+
+; GCN-LABEL: {{^}}test_s_dcache_inv_vol:
+; GCN-NEXT: ; BB#0:
+; CI-NEXT: s_dcache_inv_vol ; encoding: [0x00,0x00,0x40,0xc7]
+; VI-NEXT: s_dcache_inv_vol ; encoding: [0x00,0x00,0x88,0xc0,0x00,0x00,0x00,0x00]
+; GCN-NEXT: s_endpgm
+define void @test_s_dcache_inv_vol() #0 {
+  call void @llvm.amdgcn.s.dcache.inv.vol()
+  ret void
+}
+
+; GCN-LABEL: {{^}}test_s_dcache_inv_vol_insert_wait:
+; GCN-NEXT: ; BB#0:
+; GCN-NEXT: s_dcache_inv_vol
+; GCN-NEXT: s_waitcnt lgkmcnt(0) ; encoding
+define void @test_s_dcache_inv_vol_insert_wait() #0 {
+  call void @llvm.amdgcn.s.dcache.inv.vol()
+  br label %end
+
+end:
+  store volatile i32 3, i32 addrspace(1)* undef
+  ret void
+}
+
+attributes #0 = { nounwind }

test/CodeGen/AMDGPU/llvm.amdgcn.s.dcache.wb.ll

+; RUN: llc -march=amdgcn -mcpu=fiji -show-mc-encoding < %s | FileCheck -check-prefix=VI %s
+
+declare void @llvm.amdgcn.s.dcache.wb() #0
+
+; VI-LABEL: {{^}}test_s_dcache_wb:
+; VI-NEXT: ; BB#0:
+; VI-NEXT: s_dcache_wb ; encoding: [0x00,0x00,0x84,0xc0,0x00,0x00,0x00,0x00]
+; VI-NEXT: s_endpgm
+define void @test_s_dcache_wb() #0 {
+  call void @llvm.amdgcn.s.dcache.wb()
+  ret void
+}
+
+; VI-LABEL: {{^}}test_s_dcache_wb_insert_wait:
+; VI-NEXT: ; BB#0:
+; VI-NEXT: s_dcache_wb
+; VI-NEXT: s_waitcnt lgkmcnt(0) ; encoding
+define void @test_s_dcache_wb_insert_wait() #0 {
+  call void @llvm.amdgcn.s.dcache.wb()
+  br label %end
+
+end:
+  store volatile i32 3, i32 addrspace(1)* undef
+  ret void
+}
+
+attributes #0 = { nounwind }

test/CodeGen/AMDGPU/llvm.amdgcn.s.dcache.wb.vol.ll

+; RUN: llc -march=amdgcn -mcpu=fiji -show-mc-encoding < %s | FileCheck -check-prefix=VI %s
+
+declare void @llvm.amdgcn.s.dcache.wb.vol() #0
+
+; VI-LABEL: {{^}}test_s_dcache_wb_vol:
+; VI-NEXT: ; BB#0:
+; VI-NEXT: s_dcache_wb_vol ; encoding: [0x00,0x00,0x8c,0xc0,0x00,0x00,0x00,0x00]
+; VI-NEXT: s_endpgm
+define void @test_s_dcache_wb_vol() #0 {
+  call void @llvm.amdgcn.s.dcache.wb.vol()
+  ret void
+}
+
+; VI-LABEL: {{^}}test_s_dcache_wb_vol_insert_wait:
+; VI-NEXT: ; BB#0:
+; VI-NEXT: s_dcache_wb_vol
+; VI-NEXT: s_waitcnt lgkmcnt(0) ; encoding
+define void @test_s_dcache_wb_vol_insert_wait() #0 {
+  call void @llvm.amdgcn.s.dcache.wb.vol()
+  br label %end
+
+end:
+  store volatile i32 3, i32 addrspace(1)* undef
+  ret void
+}
+
+attributes #0 = { nounwind }

test/MC/AMDGPU/smem.s

+// RUN: llvm-mc -arch=amdgcn -mcpu=tonga -show-encoding %s | FileCheck -check-prefix=GCN -check-prefix=VI %s
+// RUN: not llvm-mc -arch=amdgcn -mcpu=tahiti %s 2>&1 | FileCheck -check-prefix=NOSI %s
+// RUN: not llvm-mc -arch=amdgcn -mcpu=bonaire %s 2>&1 | FileCheck -check-prefix=NOSI %s
+
+s_dcache_wb
+; VI: s_dcache_wb  ; encoding: [0x00,0x00,0x84,0xc0,0x00,0x00,0x00,0x00]
+; NOSI: error: instruction not supported on this GPU
+
+s_dcache_wb_vol
+; VI: s_dcache_wb_vol ; encoding: [0x00,0x00,0x8c,0xc0,0x00,0x00,0x00,0x00]
+; NOSI: error: instruction not supported on this GPU

test/MC/AMDGPU/smrd.s

 s_load_dwordx8 s[8:15], s[2:3], s4
 // GCN: s_load_dwordx8 s[8:15], s[2:3], s4 ; encoding: [0x04,0x02,0xc4,0xc0]
 
 s_load_dwordx16 s[16:31], s[2:3], 1
 // GCN: s_load_dwordx16 s[16:31], s[2:3], 0x1 ; encoding: [0x01,0x03,0x08,0xc1]
 
 s_load_dwordx16 s[16:31], s[2:3], s4
 // GCN: s_load_dwordx16 s[16:31], s[2:3], s4 ; encoding: [0x04,0x02,0x08,0xc1]
+
+s_dcache_inv
+// GCN: s_dcache_inv ; encoding: [0x00,0x00,0xc0,0xc7]
+
+s_dcache_inv_vol
+// CI: s_dcache_inv_vol ; encoding: [0x00,0x00,0x40,0xc7]
+// NOSI: error: instruction not supported on this GPU