[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 :
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 {
--- /dev/null
+; 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 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 }