[PATCH] D17892: AMDGPU/SI: Implement S_GETREG Intrinsic
Changpeng Fang via llvm-commits
llvm-commits at lists.llvm.org
Tue Mar 8 11:18:22 PST 2016
cfang updated this revision to Diff 50061.
cfang added a comment.
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.
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
@@ -416,7 +416,11 @@
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 @@
>;
//===----------------------------------------------------------------------===//
+// S_GETREG_B32 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
@@ -231,6 +231,10 @@
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]>;
-------------- next part --------------
A non-text attachment was scrubbed...
Name: D17892.50061.patch
Type: text/x-patch
Size: 2320 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20160308/f28f23d4/attachment.bin>
More information about the llvm-commits
mailing list