[PATCH] D17892: AMDGPU/SI: Implement S_GETREG Intrinsic

[Changpeng Fang via llvm-commits]

cfang created this revision.
cfang added reviewers: arsenm, tstellarAMD, scchan.
cfang added subscribers: llvm-commits, arsenm.

S_GETREG is to read hardware registers.

http://reviews.llvm.org/D17892

Files:
  include/llvm/IR/IntrinsicsAMDGPU.td
  lib/Target/AMDGPU/SIInstructions.td
  test/CodeGen/AMDGPU/llvm.amdgcn.s.getreg.ll

Index: test/CodeGen/AMDGPU/llvm.amdgcn.s.getreg.ll
===================================================================
--- /dev/null
+++ test/CodeGen/AMDGPU/llvm.amdgcn.s.getreg.ll
@@ -0,0 +1,14 @@
+; 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}
Index: lib/Target/AMDGPU/SIInstructions.td
===================================================================
--- lib/Target/AMDGPU/SIInstructions.td
+++ lib/Target/AMDGPU/SIInstructions.td
@@ -2069,6 +2069,14 @@
 >;
 
 //===----------------------------------------------------------------------===//
+// S_GETREG Intrinsic Pattern
+//===----------------------------------------------------------------------===//
+def : Pat <
+  (int_amdgcn_s_getreg imm:$simm16),
+  (S_GETREG_B32 (as_i16imm $simm16))
+>;
+
+//===----------------------------------------------------------------------===//
 // SMRD Patterns
 //===----------------------------------------------------------------------===//
 
Index: include/llvm/IR/IntrinsicsAMDGPU.td
===================================================================
--- include/llvm/IR/IntrinsicsAMDGPU.td
+++ include/llvm/IR/IntrinsicsAMDGPU.td
@@ -216,6 +216,10 @@
   GCCBuiltin<"__builtin_amdgcn_mbcnt_hi">,
   Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem]>;
 
+// llvm.amdgcn.s.getreg <simm16>
+def int_amdgcn_s_getreg :
+  Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrNoMem]>;
+
 //===----------------------------------------------------------------------===//
 // CI+ Intrinsics
 //===----------------------------------------------------------------------===//


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D17892.49836.patch
Type: text/x-patch
Size: 2048 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20160304/90242509/attachment.bin>