clang/AMDGPU: Require 16-bit-insts for half typed buffer format builtins

Typed buffer format load/store operations with 16-bit elements require
d16 support which was introduced in gfx8. These builtins previously had
no required features at all, so they were accepted (and then crashed the
backend) on targets without 16-bit support.

Diagnose these in Sema, parallel to the image builtins. The manual
verification here suprised me. The automatic builtin feature verification
is enforced in codegen, which seems like a layering violation which
should be fixed.

Co-Authored-By: Claude (Opus 4.8) <noreply@anthropic.com>

clang/include/clang/Basic/BuiltinsAMDGPU.td

@@ -261,13 +261,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)">;
 

clang/lib/Sema/SemaAMDGPU.cpp

@@ -185,6 +185,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:

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 amdgpu6.01-amd-amdhsa -emit-llvm -fcuda-is-device -o - %s | FileCheck %s
+// RUN: %clang_cc1 -O1 -triple amdgpu8.03-amd-amdhsa -emit-llvm -fcuda-is-device -o - %s | FileCheck %s
 
 #define __device__ __attribute__((device))
 

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 amdgpu6.01-unknown-unknown -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple amdgpu8.03-unknown-unknown -emit-llvm -o - %s | FileCheck %s
 
 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
 

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 amdgpu6.01-unknown-unknown -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple amdgpu8.03-unknown-unknown -emit-llvm -o - %s | FileCheck %s
 
 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
 

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 amdgpu6.01-unknown-unknown -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple amdgpu8.03-unknown-unknown -emit-llvm -o - %s | FileCheck %s
 
 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
 

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 amdgpu6.01-unknown-unknown -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple amdgpu8.03-unknown-unknown -emit-llvm -o - %s | FileCheck %s
 
 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
 

clang/test/Sema/builtins-amdgcn-d16-buffer-format-16bit-error.c

@@ -0,0 +1,19 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgpu7.00-unknown-unknown -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);
+}

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 amdgpu6.01 -verify %s -fcuda-is-device
+// RUN: %clang_cc1 -fsyntax-only -triple amdgpu8.03 -verify %s -fcuda-is-device
 // REQUIRES: amdgpu-registered-target
 
 #define __device__ __attribute__((device))