diff options
| author | Konstantin Zhuravlyov <kzhuravl_dev@outlook.com> | 2017-05-26 21:08:20 +0000 |
|---|---|---|
| committer | Konstantin Zhuravlyov <kzhuravl_dev@outlook.com> | 2017-05-26 21:08:20 +0000 |
| commit | 1f144a18ff82809c1aef1936460e30e05eac862f (patch) | |
| tree | af662eeab09bd625120410179bedf216774bc03e | |
| parent | e55b4124df4c69f6d7f6027c0921f0a0cf5cafd0 (diff) | |
| download | bcm5719-llvm-1f144a18ff82809c1aef1936460e30e05eac862f.tar.gz bcm5719-llvm-1f144a18ff82809c1aef1936460e30e05eac862f.zip | |
Resubmit r303861.
[AMDGPU] add __builtin_amdgcn_s_getpc
Patch by Tim Corringham
llvm-svn: 304033
| -rw-r--r-- | clang/include/clang/Basic/BuiltinsAMDGPU.def | 1 | ||||
| -rw-r--r-- | clang/test/CodeGenOpenCL/builtins-amdgcn.cl | 7 |
2 files changed, 8 insertions, 0 deletions
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def index a8ab657c379..6542acafe48 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.def +++ b/clang/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/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl index 748fe5e7120..f75620ba603 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl +++ b/clang/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 } |
