AMDGPU: Add another missing builtin

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@339395 91177308-0d34-0410-b5e6-96231b3b80d8

diff --git a/include/clang/Basic/BuiltinsAMDGPU.def b/include/clang/Basic/BuiltinsAMDGPU.def
index 73204dc747c32daaacf585b1a643c7cbd0c51946..35f3a887dee0f2f9833e9b409307313a35f0e88d 100644 (file)
--- a/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/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/test/CodeGenOpenCL/builtins-amdgcn-ci.cl b/test/CodeGenOpenCL/builtins-amdgcn-ci.cl
index 023e761433d3840ab295752b91f042513b8174d6..41275268dbb4b7331103bc04e1a6d46ca9d28c56 100644 (file)
--- a/test/CodeGenOpenCL/builtins-amdgcn-ci.cl
+++ b/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/test/SemaOpenCL/builtins-amdgcn-error-ci.cl b/test/SemaOpenCL/builtins-amdgcn-error-ci.cl
index 282b45fdb14465ddd0f8b57c09808fc17a58312a..2f656582bee90948c73d343a198f63c7700c79b0 100644 (file)
--- a/test/SemaOpenCL/builtins-amdgcn-error-ci.cl
+++ b/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}}
 }