AMDGPU/SI: Implement S_GETREG Intrinsic

[cfang]

Summary

S_GETREG is to read hardware registers.

[arsenm]

SC has some prettier syntax for the values being read, but that can be a separate patch

include/llvm/IR/IntrinsicsAMDGPU.td

221 ↗

(On Diff #49836)

This needs to be read only. In special cases the call site can be marked readnone

lib/Target/AMDGPU/SIInstructions.td

2076 ↗

(On Diff #49836)

You should be able to put this in the instruction definition

[cfang]

Update based on Matt's Comments:

1. Use "read only" attribute for s_getreg intrinsic;
2. Decided to use separate Pattern for the intrinsic temporarily.

[cfang]

update based on Matt's Comments

1. Create a new SOPK_32 multicalss, and rename the original one as SOPK_16 (so both are based on operand size);
2. Move the intrinsic s_getreg_b32 pattern inside the instruction definition;
3. Add checks for SI and CI devices in the test.

[arsenm]

LGTM

llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td

def int_amdgcn_s_memtime :
   GCCBuiltin<"__builtin_amdgcn_s_memtime">,
   Intrinsic<[llvm_i64_ty], [], []>;
 
 def int_amdgcn_s_sleep :
   GCCBuiltin<"__builtin_amdgcn_s_sleep">,
   Intrinsic<[], [llvm_i32_ty], []> {
 }
 
+def int_amdgcn_s_getreg :
+  GCCBuiltin<"__builtin_amdgcn_s_getreg">,
+  Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrReadMem]>;
+
 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],

llvm/trunk/lib/Target/AMDGPU/SIInstructions.td

let Defs = [SCC], isCommutable = 1, DisableEncoding = "$src0",
   defm S_ADDK_I32 : SOPK_32TIE <sopk<0x0f, 0x0e>, "s_addk_i32", []>;
   defm S_MULK_I32 : SOPK_32TIE <sopk<0x10, 0x0f>, "s_mulk_i32", []>;
 }
 
 defm S_CBRANCH_I_FORK : SOPK_m <
   sopk<0x11, 0x10>, "s_cbranch_i_fork", (outs),
   (ins SReg_64:$sdst, u16imm:$simm16), " $sdst, $simm16"
 >;
+
+let mayLoad = 1 in {
 defm S_GETREG_B32 : SOPK_32 <sopk<0x12, 0x11>, "s_getreg_b32", []>;
+}
+
 defm S_SETREG_B32 : SOPK_m <
   sopk<0x13, 0x12>, "s_setreg_b32", (outs),
   (ins SReg_32:$sdst, u16imm:$simm16), " $sdst, $simm16"
 >;
 // FIXME: Not on SI?
 //defm S_GETREG_REGRD_B32 : SOPK_32 <sopk<0x14, 0x13>, "s_getreg_regrd_b32", []>;
 defm S_SETREG_IMM32_B32 : SOPK_IMM32 <
   sopk<0x15, 0x14>, "s_setreg_imm32_b32", (outs),

 def : Pat <
   (int_SI_export imm:$en, imm:$vm, imm:$done, imm:$tgt, imm:$compr,
                  f32:$src0, f32:$src1, f32:$src2, f32:$src3),
   (EXP imm:$en, imm:$tgt, imm:$compr, imm:$done, imm:$vm,
        $src0, $src1, $src2, $src3)
 >;
 
 //===----------------------------------------------------------------------===//
+// S_GETREG_B32 Intrinsic Pattern.
+//===----------------------------------------------------------------------===//
+def : Pat <
+  (int_amdgcn_s_getreg imm:$simm16),
+  (S_GETREG_B32 (as_i16imm $simm16))
+>;
+
+//===----------------------------------------------------------------------===//
 // SMRD Patterns
 //===----------------------------------------------------------------------===//
 
 multiclass SMRD_Pattern <string Instr, ValueType vt> {
 
   // 1. IMM offset
   def : Pat <
     (smrd_load (SMRDImm i64:$sbase, i32:$offset)),

llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.s.getreg.ll

+; RUN: llc -march=amdgcn -mcpu=tahiti -verify-machineinstrs < %s | FileCheck %s
+; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=kaveri -verify-machineinstrs < %s | FileCheck %s
+; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=fiji -verify-machineinstrs < %s | FileCheck %s
+
+; FUNC-LABEL: {{^}}s_getreg_test:
+; CHECK: s_getreg_b32 s{{[0-9]+}}, 0xb206
+define void @s_getreg_test(i32 addrspace(1)* %out) { ; simm16=45574 for lds size.
+  %lds_size_64dwords = call i32 @llvm.amdgcn.s.getreg(i32 45574) #0
+  %lds_size_bytes = shl i32 %lds_size_64dwords, 8
+  store i32 %lds_size_bytes, i32 addrspace(1)* %out
+  ret void
+}
+
+declare i32 @llvm.amdgcn.s.getreg(i32) #0
+
+attributes #0 = { nounwind readonly}