https://github.com/arsenm updated https://github.com/llvm/llvm-project/pull/205367
>From dba8bec50d5b639991726d99134ab5a4b5c0f493 Mon Sep 17 00:00:00 2001 From: Matt Arsenault <[email protected]> Date: Tue, 23 Jun 2026 17:36:55 +0200 Subject: [PATCH 1/2] 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) <[email protected]> --- clang/include/clang/Basic/BuiltinsAMDGPU.td | 8 ++++---- clang/lib/Sema/SemaAMDGPU.cpp | 14 ++++++++++++++ .../builtins-amdgcn-buffer-format.hip | 2 +- .../builtins-amdgcn-raw-buffer-load-format.cl | 2 +- ...builtins-amdgcn-raw-buffer-store-format.cl | 2 +- ...iltins-amdgcn-struct-buffer-load-format.cl | 2 +- ...ltins-amdgcn-struct-buffer-store-format.cl | 2 +- ...ins-amdgcn-d16-buffer-format-16bit-error.c | 19 +++++++++++++++++++ .../SemaHIP/builtins-amdgcn-buffer-format.hip | 2 +- 9 files changed, 43 insertions(+), 10 deletions(-) create mode 100644 clang/test/Sema/builtins-amdgcn-d16-buffer-format-16bit-error.c diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.td b/clang/include/clang/Basic/BuiltinsAMDGPU.td index 77b402a8a63b2..0659aea921eca 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.td +++ b/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)">; diff --git a/clang/lib/Sema/SemaAMDGPU.cpp b/clang/lib/Sema/SemaAMDGPU.cpp index 48230fa262d5c..ad84c3335dfd2 100644 --- a/clang/lib/Sema/SemaAMDGPU.cpp +++ b/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: diff --git a/clang/test/CodeGenHIP/builtins-amdgcn-buffer-format.hip b/clang/test/CodeGenHIP/builtins-amdgcn-buffer-format.hip index 772ea080cf9c5..4e2dadaff637c 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 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)) 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 9f87769c2fd5b..b7da9d0314115 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 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 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 01e235c1d7024..5c399a39bf133 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 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 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 fac5271f9ae26..0449c2c00bbb6 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 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 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 7e04974c5e44c..3575cc8d724a2 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 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 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..df5db4aa11459 --- /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 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); +} diff --git a/clang/test/SemaHIP/builtins-amdgcn-buffer-format.hip b/clang/test/SemaHIP/builtins-amdgcn-buffer-format.hip index 04f8581c5ca75..b4b89e1382205 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 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)) >From ed8cb5f4d0fbd1c2c3cbd9a391411ae6c4472668 Mon Sep 17 00:00:00 2001 From: Matt Arsenault <[email protected]> Date: Sun, 2 Aug 2026 21:30:55 +0200 Subject: [PATCH 2/2] Use default feature check in codegen instead of sema --- clang/lib/Sema/SemaAMDGPU.cpp | 14 -------------- ...ins-amdgcn-d16-buffer-format-16bit-error.c | 19 ------------------- ...ins-amdgcn-buffer-format-d16-target-err.cl | 16 ++++++++++++++++ 3 files changed, 16 insertions(+), 33 deletions(-) delete mode 100644 clang/test/Sema/builtins-amdgcn-d16-buffer-format-16bit-error.c create mode 100644 clang/test/SemaOpenCL/builtins-amdgcn-buffer-format-d16-target-err.cl diff --git a/clang/lib/Sema/SemaAMDGPU.cpp b/clang/lib/Sema/SemaAMDGPU.cpp index ad84c3335dfd2..48230fa262d5c 100644 --- a/clang/lib/Sema/SemaAMDGPU.cpp +++ b/clang/lib/Sema/SemaAMDGPU.cpp @@ -185,20 +185,6 @@ 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/Sema/builtins-amdgcn-d16-buffer-format-16bit-error.c b/clang/test/Sema/builtins-amdgcn-d16-buffer-format-16bit-error.c deleted file mode 100644 index df5db4aa11459..0000000000000 --- a/clang/test/Sema/builtins-amdgcn-d16-buffer-format-16bit-error.c +++ /dev/null @@ -1,19 +0,0 @@ -// 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); -} diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-buffer-format-d16-target-err.cl b/clang/test/SemaOpenCL/builtins-amdgcn-buffer-format-d16-target-err.cl new file mode 100644 index 0000000000000..5f42ee2aeb2a0 --- /dev/null +++ b/clang/test/SemaOpenCL/builtins-amdgcn-buffer-format-d16-target-err.cl @@ -0,0 +1,16 @@ +// RUN: %clang_cc1 -triple amdgpu7.00-unknown-unknown -S -verify -o - %s +// REQUIRES: amdgpu-registered-target + +// Half typed buffer format load/store builtins require d16 support +// (16-bit-insts), introduced in gfx8. + +#pragma OPENCL EXTENSION cl_khr_fp16 : enable + +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(rsrc, 0, 0, 0); // expected-error {{'__builtin_amdgcn_raw_buffer_load_format_v4f16' needs target feature 16-bit-insts}} + __builtin_amdgcn_raw_buffer_store_format_v4f16(v, rsrc, 0, 0, 0); // expected-error {{'__builtin_amdgcn_raw_buffer_store_format_v4f16' needs target feature 16-bit-insts}} + v = __builtin_amdgcn_struct_buffer_load_format_v4f16(rsrc, 0, 0, 0, 0); // expected-error {{'__builtin_amdgcn_struct_buffer_load_format_v4f16' needs target feature 16-bit-insts}} + __builtin_amdgcn_struct_buffer_store_format_v4f16(v, rsrc, 0, 0, 0, 0); // expected-error {{'__builtin_amdgcn_struct_buffer_store_format_v4f16' needs target feature 16-bit-insts}} +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
