AMDGPU: Implement readcyclecounter

[nhaehnle]

Looks good to me.

lib/Target/AMDGPU/SIInstructions.td

91–100

FWIW, OpenGL (ARB_shader_clock) doesn't really care.

test/CodeGen/AMDGPU/readcyclecounter.ll

8

s_memtime needs a s_waitcnt, right? Makes sense to check for that here, but not a big deal either way.

[arsenm]

test/CodeGen/AMDGPU/readcyclecounter.ll

8

That would surprise me if it does need it since it isn't actually a load. This should probably be special cased to not require a waitcnt

[arsenm]

Disable waitcnt on s_memtime

[nhaehnle]

NAK on that. I agree that it feels odd, but...

The documentation explicitly mentions that "s_memtime counts the same as an s_load_dwordx2" (in the publicly released AMD_GCN3_Instruction_Set_Architecture.pdf, page 39 of the PDF in section 4.4, where LGKM_CNT is discussed).

[arsenm]

In D15873#319488, @nhaehnle wrote:

NAK on that. I agree that it feels odd, but...

The documentation explicitly mentions that "s_memtime counts the same as an s_load_dwordx2" (in the publicly released AMD_GCN3_Instruction_Set_Architecture.pdf, page 39 of the PDF in section 4.4, where LGKM_CNT is discussed).

OK, I didn't find this in the SI/CI manual (which seems to used old names: "smrd_fetch_time counts the same as an smrd_fetch_2.")

[arsenm]

My main question now is still whether this should map to s_memtime, or match the behavior of the HSAIL clock instruction w.r.t s_memrealtime on VI.

[arsenm]

In D15873#319491, @arsenm wrote:

In D15873#319488, @nhaehnle wrote:

NAK on that. I agree that it feels odd, but...

The documentation explicitly mentions that "s_memtime counts the same as an s_load_dwordx2" (in the publicly released AMD_GCN3_Instruction_Set_Architecture.pdf, page 39 of the PDF in section 4.4, where LGKM_CNT is discussed).

OK, I didn't find this in the SI/CI manual (which seems to used old names: "smrd_fetch_time counts the same as an smrd_fetch_2.")

Another question would be where the waitcnt should be inserted, and if this should have one inserted right after it to be more precise. Right now with the attempt to reduce waits, the s_memtime ends up being the first instruction, and has a waitcnt shared with the kernel argument load in the testcase. This probably isn't what you want

[nhaehnle]

I don't have a strong opinion s_memtime vs. s_memrealtime. Probably the least surprise is to mirror HSAIL, but I don't really care.

I am against forcing there to be a waitcnt immediately after the s_memtime, because that might cause unnecessary stalls. However, we should prevent the s_memtime itself from getting reordered. GL_ARB_shader_clock explicitly states that the clock functions are code motion barriers. Isn't the intrinsic marked as having side effects?

[arsenm]

In D15873#319504, @nhaehnle wrote:

I don't have a strong opinion s_memtime vs. s_memrealtime. Probably the least surprise is to mirror HSAIL, but I don't really care.

I am against forcing there to be a waitcnt immediately after the s_memtime, because that might cause unnecessary stalls. However, we should prevent the s_memtime itself from getting reordered. GL_ARB_shader_clock explicitly states that the clock functions are code motion barriers. Isn't the intrinsic marked as having side effects?

Yes, the intrinsic already won't be reordered with other side effects / memory writes

[nhaehnle]

Ah, now I understand. The load after the s_memtime logically belongs to the store that follows it, so there is no reordering going on. Combining the waits is exactly what we want: when s_memtime is executed, the shader hardware immediately sends the corresponding request to some other hardware block (LDS? I'm not sure). The returned cycle count will be the correct one, no matter how late the corresponding s_waitcnt is scheduled.

[nhaehnle]

So with a check in readcyclecounter.ll to ensure that lgkmcnt is present after s_memtime, this revision looks good to me (don't know if somebody else wants to chime in on the s_memtime vs. s_memrealtime issue).

[arsenm]

Make it have the same behavior as hsail clock. Add separate s_memtime and s_memrealtime intrinsic.s

[• tstellarAMD]

LGTM.

[arsenm]

r262119

include/llvm/IR/IntrinsicsAMDGPU.td

 def int_amdgcn_buffer_wbinvl1 :
   GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1">,
   Intrinsic<[], [], []>;
 
 def int_amdgcn_s_dcache_inv :
   GCCBuiltin<"__builtin_amdgcn_s_dcache_inv">,
   Intrinsic<[], [], []>;
 
+def int_amdgcn_s_memtime :
+  GCCBuiltin<"__builtin_amdgcn_s_memtime">,
+  Intrinsic<[llvm_i64_ty], [], []>;
+
 def int_amdgcn_dispatch_ptr :
   GCCBuiltin<"__builtin_amdgcn_dispatch_ptr">,
   Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [], [IntrNoMem]>;
 
 // __builtin_amdgcn_interp_p1 <i>, <attr_chan>, <attr>, <m0>
 def int_amdgcn_interp_p1 :
   GCCBuiltin<"__builtin_amdgcn_interp_p1">,
   Intrinsic<[llvm_float_ty],

 def int_amdgcn_s_dcache_wb :
   GCCBuiltin<"__builtin_amdgcn_s_dcache_wb">,
   Intrinsic<[], [], []>;
 
 def int_amdgcn_s_dcache_wb_vol :
   GCCBuiltin<"__builtin_amdgcn_s_dcache_wb_vol">,
   Intrinsic<[], [], []>;
 
+def int_amdgcn_s_memrealtime :
+  GCCBuiltin<"__builtin_amdgcn_s_memrealtime">,
+  Intrinsic<[llvm_i64_ty], [], []>;
 }

lib/Target/AMDGPU/AMDGPU.td

 >;
 
 def FeatureCIInsts : SubtargetFeature<"ci-insts",
   "CIInsts",
   "true",
   "Additional intstructions for CI+"
 >;
 
+def FeatureVIInsts : SubtargetFeature<"vi-insts",
+  "VIInsts",
+  "true",
+  "Additional intstructions for VI+"
+>;
+
 //===------------------------------------------------------------===//
 // Subtarget Features (options and debugging)
 //===------------------------------------------------------------===//
 
 // Some instructions do not support denormals despite this flag. Using
 // fp32 denormals also causes instructions to run at the double
 // precision rate for the device.
 def FeatureFP32Denormals : SubtargetFeature<"fp32-denormals",
def FeatureSeaIslands : SubtargetFeatureGeneration<"SEA_ISLANDS",
   [FeatureFP64, FeatureLocalMemorySize65536,
   FeatureWavefrontSize64, FeatureGCN, FeatureFlatAddressSpace,
   FeatureGCN1Encoding, FeatureCIInsts]
 >;
 
 def FeatureVolcanicIslands : SubtargetFeatureGeneration<"VOLCANIC_ISLANDS",
   [FeatureFP64, FeatureLocalMemorySize65536,
    FeatureWavefrontSize64, FeatureFlatAddressSpace, FeatureGCN,
-   FeatureGCN3Encoding, FeatureCIInsts]
+   FeatureGCN3Encoding, FeatureCIInsts, FeatureVIInsts]
 >;
 
 //===----------------------------------------------------------------------===//
 
 def AMDGPUInstrInfo : InstrInfo {
   let guessInstructionProperties = 1;
   let noNamedPositionallyEncodedOperands = 1;
 }

lib/Target/AMDGPU/AMDGPUSubtarget.h

private:
   int LocalMemorySize;
   unsigned MaxPrivateElementSize;
   bool EnableVGPRSpilling;
   bool SGPRInitBug;
   bool IsGCN;
   bool GCN1Encoding;
   bool GCN3Encoding;
   bool CIInsts;
+  bool VIInsts;
   bool FeatureDisable;
   int LDSBankCount;
   unsigned IsaVersion;
   bool EnableHugeScratchBuffer;
   bool EnableSIScheduler;
 
   std::unique_ptr<AMDGPUFrameLowering> FrameLowering;
   std::unique_ptr<AMDGPUTargetLowering> TLInfo;

lib/Target/AMDGPU/AMDGPUSubtarget.cpp

: AMDGPUGenSubtargetInfo(TT, GPU, FS),
       FlatAddressSpace(false), FlatForGlobal(false), EnableIRStructurizer(true),
       EnablePromoteAlloca(false),
       EnableIfCvt(true), EnableLoadStoreOpt(false),
       EnableUnsafeDSOffsetFolding(false),
       EnableXNACK(false),
       WavefrontSize(0), CFALUBug(false),
       LocalMemorySize(0), MaxPrivateElementSize(0),
       EnableVGPRSpilling(false), SGPRInitBug(false), IsGCN(false),
-      GCN1Encoding(false), GCN3Encoding(false), CIInsts(false), LDSBankCount(0),
+      GCN1Encoding(false), GCN3Encoding(false), CIInsts(false), VIInsts(false),
+      LDSBankCount(0),
       IsaVersion(ISAVersion0_0_0), EnableHugeScratchBuffer(false),
       EnableSIScheduler(false), FrameLowering(nullptr),
       InstrItins(getInstrItineraryForCPU(GPU)), TargetTriple(TT) {
 
   initializeSubtargetDependencies(TT, GPU, FS);
 
   const unsigned MaxStackAlign = 64 * 16; // Maximum stack alignment (long16)
 

lib/Target/AMDGPU/SIISelLowering.cpp

SITargetLowering::SITargetLowering(TargetMachine &TM,
 
   setOperationAction(ISD::INTRINSIC_VOID, MVT::Other, Custom);
   setOperationAction(ISD::BRCOND, MVT::Other, Custom);
   setOperationAction(ISD::BR_CC, MVT::i32, Expand);
   setOperationAction(ISD::BR_CC, MVT::i64, Expand);
   setOperationAction(ISD::BR_CC, MVT::f32, Expand);
   setOperationAction(ISD::BR_CC, MVT::f64, Expand);
 
+  // On SI this is s_memtime and s_memrealtime on VI.
+  setOperationAction(ISD::READCYCLECOUNTER, MVT::i64, Legal);
+
   for (MVT VT : MVT::integer_valuetypes()) {
     if (VT == MVT::i64)
       continue;
 
     setLoadExtAction(ISD::SEXTLOAD, VT, MVT::i1, Promote);
     setLoadExtAction(ISD::SEXTLOAD, VT, MVT::i8, Legal);
     setLoadExtAction(ISD::SEXTLOAD, VT, MVT::i16, Legal);
     setLoadExtAction(ISD::SEXTLOAD, VT, MVT::i32, Expand);

lib/Target/AMDGPU/SIInstrInfo.td

multiclass SMRD_m <smrd op, string opName, bit imm, dag outs, dag ins,
 
   // 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)]>;
+multiclass SMRD_Special <smrd op, string opName, dag outs,
+                       string opStr = "",
+                       list<dag> pattern = []> {
+  let hasSideEffects = 1 in {
+    def "" : SMRD_Pseudo <opName, outs, (ins), pattern>;
 
     let sbase = 0, offset = 0 in {
       let sdst = 0 in {
-        def _si : SMRD_Real_si <op.SI, opName, 0, (outs), (ins), opName>;
+        def _si : SMRD_Real_si <op.SI, opName, 0, outs, (ins), opName#opStr>;
       }
 
       let glc = 0, sdata = 0 in {
-        def _vi : SMRD_Real_vi <op.VI, opName, 0, (outs), (ins), opName>;
-      }
-    }
-  }
+        def _vi : SMRD_Real_vi <op.VI, opName, 0, outs, (ins), opName#opStr>;
+      }
+
+multiclass SMRD_Inval <smrd op, string opName,
+                     SDPatternOperator node> {
+  let mayStore = 1 in {
+    defm : SMRD_Special<op, opName, (outs), "", [(node)]>;
+  }
 }
 
 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;
 }
 
+class SMEM_Ret <bits<8> op, string opName, SDPatternOperator node> :
+  SMRD_Real_vi<op, opName, 0, (outs SReg_64:$dst), (ins),
+  opName#" $dst", [(set i64:$dst, (node))]> {
+  let hasSideEffects = 1;
+  let mayStore = ?;
+  let mayLoad = ?;
+  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_DWORDX8 : SMRD_Helper <
   smrd<0x0b>, "s_buffer_load_dwordx8", SReg_128, SReg_256
 >;
 
 defm S_BUFFER_LOAD_DWORDX16 : SMRD_Helper <
   smrd<0x0c>, "s_buffer_load_dwordx16", SReg_128, SReg_512
 >;
 
-//def S_MEMTIME : SMRD_ <0x0000001e, "s_memtime", []>;
+let mayStore = ? in {
+// FIXME: mayStore = ? is a workaround for tablegen bug for different
+// inferred mayStore flags for the instruction pattern vs. standalone
+// Pat. Each considers the other contradictory.
+
+defm S_MEMTIME : SMRD_Special <smrd<0x1e, 0x24>, "s_memtime",
+  (outs SReg_64:$dst), " $dst", [(set i64:$dst, (int_amdgcn_s_memtime))]
+>;
+}
 
 defm S_DCACHE_INV : SMRD_Inval <smrd<0x1f, 0x20>, "s_dcache_inv",
   int_amdgcn_s_dcache_inv>;
 
 //===----------------------------------------------------------------------===//
 // SOP1 Instructions
 //===----------------------------------------------------------------------===//
 
 let isMoveImm = 1 in {
multiclass BFMPatterns <ValueType vt, InstSI BFM, InstSI MOV> {
   >;
 }
 
 defm : BFMPatterns <i32, S_BFM_B32, S_MOV_B32>;
 // FIXME: defm : BFMPatterns <i64, S_BFM_B64, S_MOV_B64>;
 
 def : BFEPattern <V_BFE_U32, S_MOV_B32>;
 
+let Predicates = [isSICI] in {
+def : Pat <
+  (i64 (readcyclecounter)),
+  (S_MEMTIME)
+>;
+}
+
 //===----------------------------------------------------------------------===//
 // Fract Patterns
 //===----------------------------------------------------------------------===//
 
 let Predicates = [isSI] in {
 
 // V_FRACT is buggy on SI, so the F32 version is never used and (x-floor(x)) is
 // used instead. However, SI doesn't have V_FLOOR_F64, so the most efficient

lib/Target/AMDGPU/VIInstructions.td

 //===----------------------------------------------------------------------===//
 
 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>;
 
+def S_MEMREALTIME : SMEM_Ret<0x25,
+  "s_memrealtime", int_amdgcn_s_memrealtime>;
+
 } // End SIAssemblerPredicate = DisableInst, SubtargetPredicate = isVI
 
 let Predicates = [isVI] in {
 
 // 1. Offset as 20bit DWORD immediate
 def : Pat <
   (SIload_constant v4i32:$sbase, IMM20bit:$offset),
   (S_BUFFER_LOAD_DWORD_IMM $sbase, (as_i32imm $offset))
 >;
 
 //===----------------------------------------------------------------------===//
-// DPP Paterns
+// DPP Patterns
 //===----------------------------------------------------------------------===//
 
 def : Pat <
   (int_amdgcn_mov_dpp i32:$src, imm:$dpp_ctrl, imm:$bound_ctrl,
                       imm:$bank_mask, imm:$row_mask),
   (V_MOV_B32_dpp $src, (as_i32imm $dpp_ctrl), (as_i1imm $bound_ctrl),
                        (as_i32imm $bank_mask), (as_i32imm $row_mask))
 >;
 
+//===----------------------------------------------------------------------===//
+// Misc Patterns
+//===----------------------------------------------------------------------===//
+
+def : Pat <
+  (i64 (readcyclecounter)),
+  (S_MEMREALTIME)
+>;
+
 } // End Predicates = [isVI]

test/CodeGen/AMDGPU/llvm.amdgcn.s.memtime.ll

+; RUN: llc -march=amdgcn -verify-machineinstrs < %s | FileCheck -check-prefix=SI -check-prefix=GCN %s
+; RUN: llc -march=amdgcn -mcpu=tonga -verify-machineinstrs < %s | FileCheck -check-prefix=VI -check-prefix=GCN %s
+
+declare i64 @llvm.amdgcn.s.memtime() #0
+
+; GCN-LABEL: {{^}}test_s_memtime:
+; GCN-DAG: s_memtime s{{\[[0-9]+:[0-9]+\]}}
+; GCN-DAG: s_load_dwordx2
+; GCN: lgkmcnt
+; GCN: buffer_store_dwordx2
+; GCN-NOT: lgkmcnt
+; GCN: s_memtime s{{\[[0-9]+:[0-9]+\]}}
+; GCN: buffer_store_dwordx2
+define void @test_s_memtime(i64 addrspace(1)* %out) #0 {
+  %cycle0 = call i64 @llvm.amdgcn.s.memtime()
+  store volatile i64 %cycle0, i64 addrspace(1)* %out
+
+  %cycle1 = call i64 @llvm.amdgcn.s.memtime()
+  store volatile i64 %cycle1, i64 addrspace(1)* %out
+  ret void
+}
+
+attributes #0 = { nounwind }

test/CodeGen/AMDGPU/llvm.amdgcn.s.realmemtime.ll

+; RUN: llc -march=amdgcn -mcpu=tonga -verify-machineinstrs < %s | FileCheck -check-prefix=VI -check-prefix=GCN %s
+
+declare i64 @llvm.amdgcn.s.realmemtime() #0
+
+; GCN-LABEL: {{^}}test_s_realmemtime:
+; GCN-DAG: s_memrealtime s{{\[[0-9]+:[0-9]+\]}}
+; GCN-DAG: s_load_dwordx2
+; GCN: lgkmcnt
+; GCN: buffer_store_dwordx2
+; GCN-NOT: lgkmcnt
+; GCN: s_memrealtime s{{\[[0-9]+:[0-9]+\]}}
+; GCN: buffer_store_dwordx2
+define void @test_s_realmemtime(i64 addrspace(1)* %out) #0 {
+  %cycle0 = call i64 @llvm.amdgcn.s.realmemtime()
+  store volatile i64 %cycle0, i64 addrspace(1)* %out
+
+  %cycle1 = call i64 @llvm.amdgcn.s.realmemtime()
+  store volatile i64 %cycle1, i64 addrspace(1)* %out
+  ret void
+}
+
+attributes #0 = { nounwind }

test/CodeGen/AMDGPU/readcyclecounter.ll

+; RUN: llc -march=amdgcn -verify-machineinstrs < %s | FileCheck -check-prefix=SI -check-prefix=GCN %s
+; RUN: llc -march=amdgcn -mcpu=tonga -verify-machineinstrs < %s | FileCheck -check-prefix=VI -check-prefix=GCN %s
+
+declare i64 @llvm.readcyclecounter() #0
+
+; GCN-LABEL: {{^}}test_readcyclecounter:
+; SI-DAG: s_memtime s{{\[[0-9]+:[0-9]+\]}}
+; VI-DAG: s_memrealtime s{{\[[0-9]+:[0-9]+\]}}
+; GCN-DAG: s_load_dwordx2
+; GCN: lgkmcnt
+; GCN: buffer_store_dwordx2
+; GCN-NOT: lgkmcnt
+; SI: s_memtime s{{\[[0-9]+:[0-9]+\]}}
+; VI: s_memrealtime s{{\[[0-9]+:[0-9]+\]}}
+; GCN: buffer_store_dwordx2
+define void @test_readcyclecounter(i64 addrspace(1)* %out) #0 {
+  %cycle0 = call i64 @llvm.readcyclecounter()
+  store volatile i64 %cycle0, i64 addrspace(1)* %out
+
+  %cycle1 = call i64 @llvm.readcyclecounter()
+  store volatile i64 %cycle1, i64 addrspace(1)* %out
+  ret void
+}
+
+attributes #0 = { nounwind }

test/MC/AMDGPU/smrd.s

 // GCN: s_load_dwordx16 s[88:103], s[2:3], s4 ; encoding: [0x04,0x02,0x2c,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
+
+s_memtime s[0:1]
+// GCN: s_memtime s[0:1] ; encoding: [0x00,0x00,0x80,0xc7]