diff --git a/include/clang/Basic/BuiltinsAMDGPU.def b/include/clang/Basic/BuiltinsAMDGPU.def
index 73204dc747..35f3a887de 100644
--- 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 023e761433..41275268db 100644
--- 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 282b45fdb1..2f656582be 100644
--- 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}}
 }