RPCS3/llvm-mirror
1
0
Fork 0
mirror of https://github.com/RPCS3/llvm-mirror.git synced 2025-02-02 18:31:54 +00:00
Code Issues Actions Packages Projects Releases Wiki Activity

Resubmit r303859 with test fixed.

Browse Source 
[AMDGPU] add intrinsic for s_getpc

Summary: The s_getpc instruction is exposed as intrinsic llvm.amdgcn.s.getpc.

Patch by Tim Corringham

llvm-svn: 304031
This commit is contained in:
Konstantin Zhuravlyov 2017-05-26 20:38:26 +00:00
parent 961a2cb79e
commit d5fb323bee
3 changed files with 28 additions and 1 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

10
include/llvm/IR/IntrinsicsAMDGPU.td
View File

@ -566,6 +566,16 @@ def int_amdgcn_s_getreg :
[IntrReadMem, IntrSpeculatable]
>;
// int_amdgcn_s_getpc is provided to allow a specific style of position
// independent code to determine the high part of its address when it is
// known (through convention) that the code and any data of interest does
// not cross a 4Gb address boundary. Use for any other purpose may not
// produce the desired results as optimizations may cause code movement,
// especially as we explicitly use IntrNoMem to allow optimizations.
def int_amdgcn_s_getpc :
GCCBuiltin<"__builtin_amdgcn_s_getpc">,
Intrinsic<[llvm_i64_ty], [], [IntrNoMem, IntrSpeculatable]>;
// __builtin_amdgcn_interp_mov <param>, <attr_chan>, <attr>, <m0>
// param values: 0 = P10, 1 = P20, 2 = P0
def int_amdgcn_interp_mov :

4
lib/Target/AMDGPU/SOPInstructions.td
View File

@ -184,7 +184,9 @@ def S_BITSET0_B32 : SOP1_32 <"s_bitset0_b32">;
def S_BITSET0_B64 : SOP1_64_32 <"s_bitset0_b64">;
def S_BITSET1_B32 : SOP1_32 <"s_bitset1_b32">;
def S_BITSET1_B64 : SOP1_64_32 <"s_bitset1_b64">;
def S_GETPC_B64 : SOP1_64_0 <"s_getpc_b64">;
def S_GETPC_B64 : SOP1_64_0 <"s_getpc_b64",
[(set i64:$sdst, (int_amdgcn_s_getpc))]
>;
let isTerminator = 1, isBarrier = 1, SchedRW = [WriteBranch] in {

15
test/CodeGen/AMDGPU/llvm.amdgcn.s.getpc.ll Normal file
View File

@ -0,0 +1,15 @@
; RUN: llc -march=amdgcn -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
declare i64 @llvm.amdgcn.s.getpc() #0
; GCN-LABEL: {{^}}test_s_getpc:
; GCN: s_load_dwordx2
; GCN-DAG: s_getpc_b64 s{{\[[0-9]+:[0-9]+\]}}
; GCN: buffer_store_dwordx2
define amdgpu_kernel void @test_s_getpc(i64 addrspace(1)* %out) #0 {
%tmp = call i64 @llvm.amdgcn.s.getpc() #1
store volatile i64 %tmp, i64 addrspace(1)* %out, align 8
ret void
}
attributes #0 = { nounwind readnone speculatable }