diff options
| -rw-r--r-- | clang/include/clang/Basic/BuiltinsAMDGPU.def | 1 | ||||
| -rw-r--r-- | clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl | 7 | ||||
| -rw-r--r-- | clang/test/SemaOpenCL/builtins-amdgcn-error-ci.cl | 3 |
3 files changed, 10 insertions, 1 deletions
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def index 73204dc747c..35f3a887dee 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.def +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def @@ -104,6 +104,7 @@ BUILTIN(__builtin_amdgcn_ds_fmaxf, "ff*3fIiIiIb", "n") // CI+ only builtins. //===----------------------------------------------------------------------===// TARGET_BUILTIN(__builtin_amdgcn_s_dcache_inv_vol, "v", "n", "ci-insts") +TARGET_BUILTIN(__builtin_amdgcn_buffer_wbinvl1_vol, "v", "n", "ci-insts") //===----------------------------------------------------------------------===// // VI+ only builtins. diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl index 023e761433d..41275268dbb 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl @@ -10,3 +10,10 @@ void test_s_dcache_inv_vol() __builtin_amdgcn_s_dcache_inv_vol(); } +// CHECK-LABEL: @test_buffer_wbinvl1_vol +// CHECK: call void @llvm.amdgcn.buffer.wbinvl1.vol() +void test_buffer_wbinvl1_vol() +{ + __builtin_amdgcn_buffer_wbinvl1_vol(); +} + diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-error-ci.cl b/clang/test/SemaOpenCL/builtins-amdgcn-error-ci.cl index 282b45fdb14..2f656582bee 100644 --- a/clang/test/SemaOpenCL/builtins-amdgcn-error-ci.cl +++ b/clang/test/SemaOpenCL/builtins-amdgcn-error-ci.cl @@ -1,7 +1,8 @@ // REQUIRES: amdgpu-registered-target // RUN: %clang_cc1 -triple amdgcn-- -target-cpu tahiti -verify -S -o - %s -void test_ci_s_dcache_inv_vol() +void test_ci_biltins() { __builtin_amdgcn_s_dcache_inv_vol(); // expected-error {{'__builtin_amdgcn_s_dcache_inv_vol' needs target feature ci-insts}} + __builtin_amdgcn_buffer_wbinvl1_vol(); // expected-error {{'__builtin_amdgcn_buffer_wbinvl1_vol' needs target feature ci-insts}} } |
