r303861 - [AMDGPU] add __builtin_amdgcn_s_getpc

[Tim Corringham via cfe-commits]

Author: timcorringham
Date: Thu May 25 09:16:11 2017
New Revision: 303861

URL: http://llvm.org/viewvc/llvm-project?rev=303861&view=rev
Log:
[AMDGPU] add __builtin_amdgcn_s_getpc

Summary: Added the builtin corresponding to the s_getpc intrinsic added in llvm D32862

Subscribers: kzhuravl, wdng, nhaehnle, yaxunl, dstuttard, tpr, t-tye

Differential Revision: https://reviews.llvm.org/D33276

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=303861&r1=303860&r2=303861&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def (original)
+++ cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def Thu May 25 09:16:11 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=303861&r1=303860&r2=303861&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn.cl Thu May 25 09:16:11 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 }