clang/AMDGPU: Require 16-bit-insts for half typed buffer format builtins by arsenm · Pull Request #205367 · llvm/llvm-project

arsenm · 2026-06-23T15:59:19Z

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

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>

arsenm · 2026-06-23T15:59:33Z

This stack of pull requests is managed by Graphite. Learn more about stacking.

llvmorg-github-actions · 2026-06-23T16:00:09Z

@llvm/pr-subscribers-clang

@llvm/pr-subscribers-backend-amdgpu

Author: Matt Arsenault (arsenm)

Changes

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>

Full diff: https://github.com/llvm/llvm-project/pull/205367.diff

9 Files Affected:

(modified) clang/include/clang/Basic/BuiltinsAMDGPU.td (+4-4)
(modified) clang/lib/Sema/SemaAMDGPU.cpp (+14)
(modified) clang/test/CodeGenHIP/builtins-amdgcn-buffer-format.hip (+1-1)
(modified) clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-load-format.cl (+1-1)
(modified) clang/test/CodeGenOpenCL/builtins-amdgcn-raw-buffer-store-format.cl (+1-1)
(modified) clang/test/CodeGenOpenCL/builtins-amdgcn-struct-buffer-load-format.cl (+1-1)
(modified) clang/test/CodeGenOpenCL/builtins-amdgcn-struct-buffer-store-format.cl (+1-1)
(added) clang/test/Sema/builtins-amdgcn-d16-buffer-format-16bit-error.c (+19)
(modified) clang/test/SemaHIP/builtins-amdgcn-buffer-format.hip (+1-1)

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))

ranapratap55

LGTM.

shiltian

We definitely don't need to do the feature check manually.

arsenm · 2026-06-27T16:13:37Z

We definitely don't need to do the feature check manually.

You do to get the check performed in the correct place in Sema. I would consider relying on the default error in codegen to be a regression and moving the default error a separate problem

shiltian · 2026-06-27T18:55:35Z

The builtin infrastructure can perform this check automatically, and that's how all the other builtins handle it today. If you don't think this is the right place for the check, I'd rather move the generic feature check somewhere else.

Also, this check doesn't work, as shown by ROCM-26503. You need more than this to handle the case correctly.

arsenm mentioned this pull request Jun 23, 2026

clang/AMDGPU: Require 16-bit-insts for half typed image builtins #205368

Open

arsenm added the backend:AMDGPU label Jun 23, 2026 — with Graphite App

arsenm requested review from ranapratap55 and shiltian June 23, 2026 15:59

arsenm marked this pull request as ready for review June 23, 2026 16:00

llvmorg-github-actions Bot added clang Clang issues not falling into any other category clang:frontend Language frontend issues, e.g. anything involving "Sema" labels Jun 23, 2026

ranapratap55 approved these changes Jun 27, 2026

View reviewed changes

shiltian requested changes Jun 27, 2026

View reviewed changes

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

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

clang/AMDGPU: Require 16-bit-insts for half typed buffer format builtins#205367
arsenm wants to merge 1 commit into
mainfrom
users/arsenm/clang/amdgpu-fix-accepting-d16-format-builtins-gfx6-gfx7

arsenm commented Jun 23, 2026

Uh oh!

arsenm commented Jun 23, 2026

Uh oh!

llvmorg-github-actions Bot commented Jun 23, 2026 •

edited

Loading

Uh oh!

ranapratap55 left a comment

Uh oh!

shiltian left a comment

Uh oh!

arsenm commented Jun 27, 2026

Uh oh!

shiltian commented Jun 27, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

Uh oh!

Conversation

arsenm commented Jun 23, 2026

Uh oh!

arsenm commented Jun 23, 2026

Uh oh!

llvmorg-github-actions Bot commented Jun 23, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

ranapratap55 left a comment

Choose a reason for hiding this comment

Uh oh!

shiltian left a comment

Choose a reason for hiding this comment

Uh oh!

arsenm commented Jun 27, 2026

Uh oh!

shiltian commented Jun 27, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

llvmorg-github-actions Bot commented Jun 23, 2026 •

edited

Loading