From 40ed16ee38828a60d1eb11122170d8ac8611422c Mon Sep 17 00:00:00 2001 From: Tim Corringham Date: Thu, 25 May 2017 14:16:11 +0000 Subject: [PATCH] [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 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@303861 91177308-0d34-0410-b5e6-96231b3b80d8 --- include/clang/Basic/BuiltinsAMDGPU.def | 1 + test/CodeGenOpenCL/builtins-amdgcn.cl | 7 +++++++ 2 files changed, 8 insertions(+) diff --git a/include/clang/Basic/BuiltinsAMDGPU.def b/include/clang/Basic/BuiltinsAMDGPU.def index a8ab657c37..6542acafe4 100644 --- a/include/clang/Basic/BuiltinsAMDGPU.def +++ b/include/clang/Basic/BuiltinsAMDGPU.def @@ -36,6 +36,7 @@ BUILTIN(__builtin_amdgcn_workitem_id_z, "Ui", "nc") // 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") diff --git a/test/CodeGenOpenCL/builtins-amdgcn.cl b/test/CodeGenOpenCL/builtins-amdgcn.cl index 748fe5e712..f75620ba60 100644 --- a/test/CodeGenOpenCL/builtins-amdgcn.cl +++ b/test/CodeGenOpenCL/builtins-amdgcn.cl @@ -481,6 +481,13 @@ void test_fmed3_f32(global float* out, float a, float b, float c) *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 } -- 2.40.0