diff options
| author | Matt Arsenault <Matthew.Arsenault@amd.com> | 2018-08-09 22:18:37 +0000 |
|---|---|---|
| committer | Matt Arsenault <Matthew.Arsenault@amd.com> | 2018-08-09 22:18:37 +0000 |
| commit | 94abc57e3751f591fdb0655419eb129252d1f9c7 (patch) | |
| tree | c3a89b3cc4fb0efd8787a46353b780240ba63b1c | |
| parent | 941095100fd8baf6c9d1dc8300d7d8637cb311bb (diff) | |
| download | bcm5719-llvm-94abc57e3751f591fdb0655419eb129252d1f9c7.tar.gz bcm5719-llvm-94abc57e3751f591fdb0655419eb129252d1f9c7.zip | |
AMDGPU: Add another missing builtin
llvm-svn: 339395
| -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}} } |
