diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.td b/clang/include/clang/Basic/BuiltinsAMDGPU.td
index ccbf2f97a1313..6c4b135c6077b 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.td
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.td
@@ -257,13 +257,13 @@ def __builtin_amdgcn_raw_buffer_load_b96 : AMDGPUBuiltin<"_ExtVector<3, unsigned
 def __builtin_amdgcn_raw_buffer_load_b128 : AMDGPUBuiltin<"_ExtVector<4, unsigned int>(__amdgpu_buffer_rsrc_t, int, int, _Constant int)">;
 
 def __builtin_amdgcn_raw_buffer_load_format_v4f32 : AMDGPUBuiltin<"_ExtVector<4, float>(__amdgpu_buffer_rsrc_t, int, int, _Constant int)">;
-def __builtin_amdgcn_raw_buffer_load_format_v4f16 : AMDGPUBuiltin<"_ExtVector<4, _Float16>(__amdgpu_buffer_rsrc_t, int, int, _Constant int)">;
+def __builtin_amdgcn_raw_buffer_load_format_v4f16 : AMDGPUBuiltin<"_ExtVector<4, _Float16>(__amdgpu_buffer_rsrc_t, int, int, _Constant int)", [], "16-bit-insts">;
 def __builtin_amdgcn_raw_buffer_store_format_v4f32 : AMDGPUBuiltin<"void(_ExtVector<4, float>, __amdgpu_buffer_rsrc_t, int, int, _Constant int)">;
-def __builtin_amdgcn_raw_buffer_store_format_v4f16 : AMDGPUBuiltin<"void(_ExtVector<4, _Float16>, __amdgpu_buffer_rsrc_t, int, int, _Constant int)">;
+def __builtin_amdgcn_raw_buffer_store_format_v4f16 : AMDGPUBuiltin<"void(_ExtVector<4, _Float16>, __amdgpu_buffer_rsrc_t, int, int, _Constant int)", [], "16-bit-insts">;
 def __builtin_amdgcn_struct_buffer_load_format_v4f32 : AMDGPUBuiltin<"_ExtVector<4, float>(__amdgpu_buffer_rsrc_t, int, int, int, _Constant int)">;
-def __builtin_amdgcn_struct_buffer_load_format_v4f16 : AMDGPUBuiltin<"_ExtVector<4, _Float16>(__amdgpu_buffer_rsrc_t, int, int, int, _Constant int)">;
+def __builtin_amdgcn_struct_buffer_load_format_v4f16 : AMDGPUBuiltin<"_ExtVector<4, _Float16>(__amdgpu_buffer_rsrc_t, int, int, int, _Constant int)", [], "16-bit-insts">;
 def __builtin_amdgcn_struct_buffer_store_format_v4f32 : AMDGPUBuiltin<"void(_ExtVector<4, float>, __amdgpu_buffer_rsrc_t, int, int, int, _Constant int)">;
-def __builtin_amdgcn_struct_buffer_store_format_v4f16 : AMDGPUBuiltin<"void(_ExtVector<4, _Float16>, __amdgpu_buffer_rsrc_t, int, int, int, _Constant int)">;
+def __builtin_amdgcn_struct_buffer_store_format_v4f16 : AMDGPUBuiltin<"void(_ExtVector<4, _Float16>, __amdgpu_buffer_rsrc_t, int, int, int, _Constant int)", [], "16-bit-insts">;
 
 def __builtin_amdgcn_raw_ptr_buffer_atomic_add_i32 : AMDGPUBuiltin<"int(int, __amdgpu_buffer_rsrc_t, int, int, _Constant int)">;
 
diff --git a/clang/lib/Sema/SemaAMDGPU.cpp b/clang/lib/Sema/SemaAMDGPU.cpp
index 757cdfbf20819..975dd2efb729c 100644
--- a/clang/lib/Sema/SemaAMDGPU.cpp
+++ b/clang/lib/Sema/SemaAMDGPU.cpp
@@ -168,6 +168,20 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,
   case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b64:
   case AMDGPU::BI__builtin_amdgcn_global_load_monitor_b128:
     return checkAtomicMonitorLoad(TheCall);
+  case AMDGPU::BI__builtin_amdgcn_raw_buffer_load_format_v4f16:
+  case AMDGPU::BI__builtin_amdgcn_raw_buffer_store_format_v4f16:
+  case AMDGPU::BI__builtin_amdgcn_struct_buffer_load_format_v4f16:
+  case AMDGPU::BI__builtin_amdgcn_struct_buffer_store_format_v4f16: {
+    StringRef FeatureList(
+        getASTContext().BuiltinInfo.getRequiredFeatures(BuiltinID));
+    if (!Builtin::evaluateRequiredTargetFeatures(FeatureList,
+                                                 CallerFeatureMap)) {
+      Diag(TheCall->getBeginLoc(), diag::err_builtin_needs_feature)
+          << FD->getDeclName() << FeatureList;
+      return false;
+    }
+    return false;
+  }
   case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f32_i32:
   case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f32_i32:
   case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f16_i32:
diff --git a/clang/test/CodeGenHIP/builtins-amdgcn-buffer-format.hip b/clang/test/CodeGenHIP/builtins-amdgcn-buffer-format.hip
index 603e6522cd38c..59165eba3a077 100644
--- a/clang/test/CodeGenHIP/builtins-amdgcn-buffer-format.hip
+++ b/clang/test/CodeGenHIP/builtins-amdgcn-buffer-format.hip
@@ -1,6 +1,6 @@
 // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
 // REQUIRES: amdgpu-registered-target
-// RUN: %clang_cc1 -O1 -triple amdgcn-amd-amdhsa -target-cpu verde -emit-llvm -fcuda-is-device -o - %s | FileCheck %s
+// RUN: %clang_cc1 -O1 -triple amdgcn-amd-amdhsa -target-cpu gfx803 -emit-llvm -fcuda-is-device -o - %s | FileCheck %s
 
 #define __device__ __attribute__((device))
 
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-load-format.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-load-format.cl
index 5c2e3e1a24862..b031a9a0c84f0 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-load-format.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-load-format.cl
@@ -1,6 +1,6 @@
 // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
 // REQUIRES: amdgpu-registered-target
-// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu verde -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx803 -emit-llvm -o - %s | FileCheck %s
 
 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
 
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-store-format.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-store-format.cl
index b10c6d59635f4..5c803bd1b8397 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-store-format.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-store-format.cl
@@ -1,6 +1,6 @@
 // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
 // REQUIRES: amdgpu-registered-target
-// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu verde -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx803 -emit-llvm -o - %s | FileCheck %s
 
 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
 
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-struct-buffer-load-format.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-struct-buffer-load-format.cl
index c31c6ed82b82f..b2d6536c75e34 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-struct-buffer-load-format.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-struct-buffer-load-format.cl
@@ -1,6 +1,6 @@
 // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
 // REQUIRES: amdgpu-registered-target
-// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu verde -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx803 -emit-llvm -o - %s | FileCheck %s
 
 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
 
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-struct-buffer-store-format.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-struct-buffer-store-format.cl
index b30a46eb78f32..110aa2cac9ecd 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-struct-buffer-store-format.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-struct-buffer-store-format.cl
@@ -1,6 +1,6 @@
 // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
 // REQUIRES: amdgpu-registered-target
-// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu verde -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx803 -emit-llvm -o - %s | FileCheck %s
 
 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
 
diff --git a/clang/test/Sema/builtins-amdgcn-d16-buffer-format-16bit-error.c b/clang/test/Sema/builtins-amdgcn-d16-buffer-format-16bit-error.c
new file mode 100644
index 0000000000000..f9ddb24910f12
--- /dev/null
+++ b/clang/test/Sema/builtins-amdgcn-d16-buffer-format-16bit-error.c
@@ -0,0 +1,19 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx700 -verify -fsyntax-only %s
+
+// Verify that half typed buffer format load/store intrinsics require
+// 16-bit-insts.
+
+typedef _Float16 half;
+typedef half half4 __attribute__((ext_vector_type(4)));
+
+void test(half4 v, __amdgpu_buffer_rsrc_t rsrc) {
+  v = __builtin_amdgcn_raw_buffer_load_format_v4f16( // expected-error {{needs target feature 16-bit-insts}}
+      rsrc, 0, 0, 0);
+  __builtin_amdgcn_raw_buffer_store_format_v4f16( // expected-error {{needs target feature 16-bit-insts}}
+      v, rsrc, 0, 0, 0);
+  v = __builtin_amdgcn_struct_buffer_load_format_v4f16( // expected-error {{needs target feature 16-bit-insts}}
+      rsrc, 0, 0, 0, 0);
+  __builtin_amdgcn_struct_buffer_store_format_v4f16( // expected-error {{needs target feature 16-bit-insts}}
+      v, rsrc, 0, 0, 0, 0);
+}
diff --git a/clang/test/SemaHIP/builtins-amdgcn-buffer-format.hip b/clang/test/SemaHIP/builtins-amdgcn-buffer-format.hip
index 15f02f821b0ba..76e1946e824ba 100644
--- a/clang/test/SemaHIP/builtins-amdgcn-buffer-format.hip
+++ b/clang/test/SemaHIP/builtins-amdgcn-buffer-format.hip
@@ -1,5 +1,5 @@
 // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6
-// RUN: %clang_cc1 -fsyntax-only -triple amdgcn -target-cpu verde -verify %s -fcuda-is-device
+// RUN: %clang_cc1 -fsyntax-only -triple amdgcn -target-cpu gfx803 -verify %s -fcuda-is-device
 // REQUIRES: amdgpu-registered-target
 
 #define __device__ __attribute__((device))