Index: include/clang/Basic/BuiltinsAMDGPU.def
===================================================================
--- include/clang/Basic/BuiltinsAMDGPU.def
+++ include/clang/Basic/BuiltinsAMDGPU.def
@@ -104,6 +104,7 @@
 // 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.
Index: test/CodeGenOpenCL/builtins-amdgcn-ci.cl
===================================================================
--- test/CodeGenOpenCL/builtins-amdgcn-ci.cl
+++ test/CodeGenOpenCL/builtins-amdgcn-ci.cl
@@ -10,3 +10,10 @@
   __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();
+}
+
Index: test/SemaOpenCL/builtins-amdgcn-error-ci.cl
===================================================================
--- test/SemaOpenCL/builtins-amdgcn-error-ci.cl
+++ 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}}
 }