r304033 - Resubmit r303861.

[Konstantin Zhuravlyov via cfe-commits]

Author: kzhuravl
Date: Fri May 26 16:08:20 2017
New Revision: 304033

URL: http://llvm.org/viewvc/llvm-project?rev=304033&view=rev
Log:
Resubmit r303861.

[AMDGPU] add __builtin_amdgcn_s_getpc

Patch by Tim Corringham

Modified:
    cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
    cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl

Modified: cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def?rev=304033&r1=304032&r2=304033&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def (original)
+++ cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def Fri May 26 16:08:20 2017
@@ -36,6 +36,7 @@ BUILTIN(__builtin_amdgcn_workitem_id_z,
 // Instruction builtins.
 //===----------------------------------------------------------------------===//
 BUILTIN(__builtin_amdgcn_s_getreg, "UiIi", "n")
+BUILTIN(__builtin_amdgcn_s_getpc, "LUi", "n")
 BUILTIN(__builtin_amdgcn_s_waitcnt, "vIi", "n")
 BUILTIN(__builtin_amdgcn_s_sendmsg, "vIiUi", "n")
 BUILTIN(__builtin_amdgcn_s_sendmsghalt, "vIiUi", "n")

Modified: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl?rev=304033&r1=304032&r2=304033&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl Fri May 26 16:08:20 2017
@@ -481,6 +481,13 @@ void test_fmed3_f32(global float* out, f
   *out = __builtin_amdgcn_fmed3f(a, b, c);
 }
 
+// CHECK-LABEL: @test_s_getpc
+// CHECK: call i64 @llvm.amdgcn.s.getpc()
+void test_s_getpc(global ulong* out)
+{
+  *out = __builtin_amdgcn_s_getpc();
+}
+
 // CHECK-DAG: [[WI_RANGE]] = !{i32 0, i32 1024}
 // CHECK-DAG: attributes #[[NOUNWIND_READONLY:[0-9]+]] = { nounwind readonly }
 // CHECK-DAG: attributes #[[READ_EXEC_ATTRS]] = { convergent }