Index: include/llvm/IR/IntrinsicsAMDGPU.td
===================================================================
--- include/llvm/IR/IntrinsicsAMDGPU.td
+++ include/llvm/IR/IntrinsicsAMDGPU.td
@@ -566,6 +566,14 @@
   [IntrReadMem, IntrSpeculatable]
 >;
 
+// Note that if a precise address is required then a barrier is needed to
+// prevent optimizations moving the s_getpc instruction, as not even
+// IntrWriteMem is sufficient to prevent all movement. We use IntrNoMem
+// here to allow optimizations when a precise address isn't required.
+def int_amdgcn_s_getpc :
+  GCCBuiltin<"__builtin_amdgcn_s_getpc">,
+  Intrinsic<[llvm_i64_ty], [], [IntrNoMem]>;
+
 // __builtin_amdgcn_interp_mov <param>, <attr_chan>, <attr>, <m0>
 // param values: 0 = P10, 1 = P20, 2 = P0
 def int_amdgcn_interp_mov :
Index: lib/Target/AMDGPU/SOPInstructions.td
===================================================================
--- lib/Target/AMDGPU/SOPInstructions.td
+++ lib/Target/AMDGPU/SOPInstructions.td
@@ -931,6 +931,14 @@
 let Predicates = [isGCN] in {
 
 //===----------------------------------------------------------------------===//
+// S_GETPC_B64 Intrinsic Pattern.
+//===----------------------------------------------------------------------===//
+def : Pat <
+  (int_amdgcn_s_getpc),
+  (S_GETPC_B64)
+>;
+
+//===----------------------------------------------------------------------===//
 // S_GETREG_B32 Intrinsic Pattern.
 //===----------------------------------------------------------------------===//
 def : Pat <
Index: test/CodeGen/AMDGPU/llvm.amdgcn.s.getpc.ll
===================================================================
--- /dev/null
+++ test/CodeGen/AMDGPU/llvm.amdgcn.s.getpc.ll
@@ -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 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 4
+  ret void
+}
+
+attributes #0 = { nounwind }