[llvm] r263124 - AMDGPU/SI: Define S_GETREG Intrinsic

[Changpeng Fang via llvm-commits]

Author: chfang
Date: Thu Mar 10 10:47:15 2016
New Revision: 263124

URL: http://llvm.org/viewvc/llvm-project?rev=263124&view=rev
Log:
AMDGPU/SI: Define S_GETREG Intrinsic

Summary:
 Define s_getreg intrinsic to generate s_getreg instruction to read
hardware registers.

Reviewers: tstellarAMD, arsenm

Subscribers: llvm-commits, arsenm

Differential Revision: http://reviews.llvm.org/D17892

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

Modified: llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td?rev=263124&r1=263123&r2=263124&view=diff
==============================================================================
--- llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td (original)
+++ llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td Thu Mar 10 10:47:15 2016
@@ -231,6 +231,10 @@ def int_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]>;

Modified: llvm/trunk/lib/Target/AMDGPU/SIInstructions.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/SIInstructions.td?rev=263124&r1=263123&r2=263124&view=diff
==============================================================================
--- llvm/trunk/lib/Target/AMDGPU/SIInstructions.td (original)
+++ llvm/trunk/lib/Target/AMDGPU/SIInstructions.td Thu Mar 10 10:47:15 2016
@@ -416,7 +416,11 @@ 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"
@@ -2090,6 +2094,14 @@ def : Pat <
 >;
 
 //===----------------------------------------------------------------------===//
+// S_GETREG_B32 Intrinsic Pattern.
+//===----------------------------------------------------------------------===//
+def : Pat <
+  (int_amdgcn_s_getreg imm:$simm16),
+  (S_GETREG_B32 (as_i16imm $simm16))
+>;
+
+//===----------------------------------------------------------------------===//
 // SMRD Patterns
 //===----------------------------------------------------------------------===//
 

Added: llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.s.getreg.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.s.getreg.ll?rev=263124&view=auto
==============================================================================
--- llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.s.getreg.ll (added)
+++ llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.s.getreg.ll Thu Mar 10 10:47:15 2016
@@ -0,0 +1,16 @@
+; 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}