diff options
| author | Jan Vesely <jan.vesely@rutgers.edu> | 2017-02-25 04:20:20 +0000 |
|---|---|---|
| committer | Jan Vesely <jan.vesely@rutgers.edu> | 2017-02-25 04:20:20 +0000 |
| commit | b2e0f431f035760dc111ca3d93ae4b9e28403dca (patch) | |
| tree | 572f9fc117f1bfcf9d2666121f45a61e0e58b4bd | |
| parent | f94f3debabb22f14da326f26d7307d592f6b3a47 (diff) | |
AMDGPU: export s_waitcnt builtin
Differential Revision: https://reviews.llvm.org/D30359
git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@296239 91177308-0d34-0410-b5e6-96231b3b80d8
| -rw-r--r-- | include/clang/Basic/BuiltinsAMDGPU.def | 1 | ||||
| -rw-r--r-- | test/CodeGenOpenCL/builtins-amdgcn.cl | 7 | ||||
| -rw-r--r-- | test/SemaOpenCL/builtins-amdgcn-error.cl | 5 |
3 files changed, 13 insertions, 0 deletions
diff --git a/include/clang/Basic/BuiltinsAMDGPU.def b/include/clang/Basic/BuiltinsAMDGPU.def index 48080d71aa..f39a24157a 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_waitcnt, "vIi", "n") BUILTIN(__builtin_amdgcn_s_barrier, "v", "n") BUILTIN(__builtin_amdgcn_wave_barrier, "v", "n") BUILTIN(__builtin_amdgcn_div_scale, "dddbb*", "n") diff --git a/test/CodeGenOpenCL/builtins-amdgcn.cl b/test/CodeGenOpenCL/builtins-amdgcn.cl index 0f6eaef50b..e9e765b521 100644 --- a/test/CodeGenOpenCL/builtins-amdgcn.cl +++ b/test/CodeGenOpenCL/builtins-amdgcn.cl @@ -263,6 +263,13 @@ void test_class_f64(global double* out, double a, int b) *out = __builtin_amdgcn_class(a, b); } +// CHECK-LABEL: @test_s_waitcnt +// CHECK: call void @llvm.amdgcn.s.waitcnt( +void test_s_waitcnt() +{ + __builtin_amdgcn_s_waitcnt(0); +} + // CHECK-LABEL: @test_s_barrier // CHECK: call void @llvm.amdgcn.s.barrier( void test_s_barrier() diff --git a/test/SemaOpenCL/builtins-amdgcn-error.cl b/test/SemaOpenCL/builtins-amdgcn-error.cl index d8576e2bb5..5e45f7e76c 100644 --- a/test/SemaOpenCL/builtins-amdgcn-error.cl +++ b/test/SemaOpenCL/builtins-amdgcn-error.cl @@ -18,6 +18,11 @@ void test_s_sleep(int x) __builtin_amdgcn_s_sleep(x); // expected-error {{argument to '__builtin_amdgcn_s_sleep' must be a constant integer}} } +void test_s_waitcnt(int x) +{ + __builtin_amdgcn_s_waitcnt(x); // expected-error {{argument to '__builtin_amdgcn_s_waitcnt' must be a constant integer}} +} + void test_s_incperflevel(int x) { __builtin_amdgcn_s_incperflevel(x); // expected-error {{argument to '__builtin_amdgcn_s_incperflevel' must be a constant integer}} |