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

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

Typed image load/store operations with 16-bit elements require d16
support which was introduced in gfx8. They were previously gated only
on image-insts, so they were wrongly accepted on targets that have
images but lack 16-bit support (e.g. gfx700), where the backend then
fails to select.

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

Typed image load/store operations with 16-bit elements require d16 support which was introduced in gfx8. They were previously gated only on image-insts, so they were wrongly accepted on targets that have images but lack 16-bit support (e.g. gfx700), where the backend then fails to select. Co-Authored-By: Claude (Opus 4.8) <noreply@anthropic.com>

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

Warning

This pull request is not mergeable via GitHub because a downstack PR is open. Once all requirements are satisfied, merge this PR as a stack on Graphite.
Learn more

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

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

@llvm/pr-subscribers-clang

@llvm/pr-subscribers-backend-amdgpu

Author: Matt Arsenault (arsenm)

Changes

Typed image load/store operations with 16-bit elements require d16
support which was introduced in gfx8. They were previously gated only
on image-insts, so they were wrongly accepted on targets that have
images but lack 16-bit support (e.g. gfx700), where the backend then
fails to select.

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

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

2 Files Affected:

(modified) clang/include/clang/Basic/BuiltinsAMDGPU.td (+24-24)
(added) clang/test/Sema/builtins-amdgcn-d16-image-16bit-error.c (+59)

diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.td b/clang/include/clang/Basic/BuiltinsAMDGPU.td
index 6c4b135c6077b..ab3d91d37a8e2 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.td
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.td
@@ -1244,61 +1244,61 @@ def __builtin_amdgcn_cooperative_atomic_store_8x16B : AMDGPUBuiltin<"void(_ExtVe
 // Image builtins
 //===----------------------------------------------------------------------===//
 def __builtin_amdgcn_image_load_1d_v4f32_i32 : AMDGPUBuiltin<"_ExtVector<4, float>(int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
-def __builtin_amdgcn_image_load_1d_v4f16_i32 : AMDGPUBuiltin<"_ExtVector<4, _Float16>(int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
+def __builtin_amdgcn_image_load_1d_v4f16_i32 : AMDGPUBuiltin<"_ExtVector<4, _Float16>(int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts,16-bit-insts">;
 def __builtin_amdgcn_image_load_1darray_v4f32_i32 : AMDGPUBuiltin<"_ExtVector<4, float>(int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
-def __builtin_amdgcn_image_load_1darray_v4f16_i32 : AMDGPUBuiltin<"_ExtVector<4, _Float16>(int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
+def __builtin_amdgcn_image_load_1darray_v4f16_i32 : AMDGPUBuiltin<"_ExtVector<4, _Float16>(int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts,16-bit-insts">;
 def __builtin_amdgcn_image_load_2d_f32_i32 : AMDGPUBuiltin<"float(int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
 def __builtin_amdgcn_image_load_2d_v4f32_i32 : AMDGPUBuiltin<"_ExtVector<4, float>(int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
-def __builtin_amdgcn_image_load_2d_v4f16_i32 : AMDGPUBuiltin<"_ExtVector<4, _Float16>(int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
+def __builtin_amdgcn_image_load_2d_v4f16_i32 : AMDGPUBuiltin<"_ExtVector<4, _Float16>(int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts,16-bit-insts">;
 def __builtin_amdgcn_image_load_2darray_f32_i32 : AMDGPUBuiltin<"float(int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
 def __builtin_amdgcn_image_load_2darray_v4f32_i32 : AMDGPUBuiltin<"_ExtVector<4, float>(int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
-def __builtin_amdgcn_image_load_2darray_v4f16_i32 : AMDGPUBuiltin<"_ExtVector<4, _Float16>(int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
+def __builtin_amdgcn_image_load_2darray_v4f16_i32 : AMDGPUBuiltin<"_ExtVector<4, _Float16>(int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts,16-bit-insts">;
 def __builtin_amdgcn_image_load_3d_v4f32_i32 : AMDGPUBuiltin<"_ExtVector<4, float>(int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
-def __builtin_amdgcn_image_load_3d_v4f16_i32 : AMDGPUBuiltin<"_ExtVector<4, _Float16>(int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
+def __builtin_amdgcn_image_load_3d_v4f16_i32 : AMDGPUBuiltin<"_ExtVector<4, _Float16>(int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts,16-bit-insts">;
 def __builtin_amdgcn_image_load_cube_v4f32_i32 : AMDGPUBuiltin<"_ExtVector<4, float>(int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
-def __builtin_amdgcn_image_load_cube_v4f16_i32 : AMDGPUBuiltin<"_ExtVector<4, _Float16>(int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
+def __builtin_amdgcn_image_load_cube_v4f16_i32 : AMDGPUBuiltin<"_ExtVector<4, _Float16>(int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts,16-bit-insts">;
 def __builtin_amdgcn_image_load_mip_1d_v4f32_i32 : AMDGPUBuiltin<"_ExtVector<4, float>(int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
-def __builtin_amdgcn_image_load_mip_1d_v4f16_i32 : AMDGPUBuiltin<"_ExtVector<4, _Float16>(int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
+def __builtin_amdgcn_image_load_mip_1d_v4f16_i32 : AMDGPUBuiltin<"_ExtVector<4, _Float16>(int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts,16-bit-insts">;
 def __builtin_amdgcn_image_load_mip_1darray_v4f32_i32 : AMDGPUBuiltin<"_ExtVector<4, float>(int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
-def __builtin_amdgcn_image_load_mip_1darray_v4f16_i32 : AMDGPUBuiltin<"_ExtVector<4, _Float16>(int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
+def __builtin_amdgcn_image_load_mip_1darray_v4f16_i32 : AMDGPUBuiltin<"_ExtVector<4, _Float16>(int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts,16-bit-insts">;
 def __builtin_amdgcn_image_load_mip_2d_f32_i32 : AMDGPUBuiltin<"float(int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
 def __builtin_amdgcn_image_load_mip_2d_v4f32_i32 : AMDGPUBuiltin<"_ExtVector<4, float>(int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
-def __builtin_amdgcn_image_load_mip_2d_v4f16_i32 : AMDGPUBuiltin<"_ExtVector<4, _Float16>(int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
+def __builtin_amdgcn_image_load_mip_2d_v4f16_i32 : AMDGPUBuiltin<"_ExtVector<4, _Float16>(int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts,16-bit-insts">;
 def __builtin_amdgcn_image_load_mip_2darray_f32_i32 : AMDGPUBuiltin<"float(int, int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
 def __builtin_amdgcn_image_load_mip_2darray_v4f32_i32 : AMDGPUBuiltin<"_ExtVector<4, float>(int, int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
-def __builtin_amdgcn_image_load_mip_2darray_v4f16_i32 : AMDGPUBuiltin<"_ExtVector<4, _Float16>(int, int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
+def __builtin_amdgcn_image_load_mip_2darray_v4f16_i32 : AMDGPUBuiltin<"_ExtVector<4, _Float16>(int, int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts,16-bit-insts">;
 def __builtin_amdgcn_image_load_mip_3d_v4f32_i32 : AMDGPUBuiltin<"_ExtVector<4, float>(int, int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
-def __builtin_amdgcn_image_load_mip_3d_v4f16_i32 : AMDGPUBuiltin<"_ExtVector<4, _Float16>(int, int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
+def __builtin_amdgcn_image_load_mip_3d_v4f16_i32 : AMDGPUBuiltin<"_ExtVector<4, _Float16>(int, int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts,16-bit-insts">;
 def __builtin_amdgcn_image_load_mip_cube_v4f32_i32 : AMDGPUBuiltin<"_ExtVector<4, float>(int, int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
-def __builtin_amdgcn_image_load_mip_cube_v4f16_i32 : AMDGPUBuiltin<"_ExtVector<4, _Float16>(int, int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
+def __builtin_amdgcn_image_load_mip_cube_v4f16_i32 : AMDGPUBuiltin<"_ExtVector<4, _Float16>(int, int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts,16-bit-insts">;
 def __builtin_amdgcn_image_store_1d_v4f32_i32 : AMDGPUBuiltin<"void(_ExtVector<4, float>, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
-def __builtin_amdgcn_image_store_1d_v4f16_i32 : AMDGPUBuiltin<"void(_ExtVector<4, _Float16>, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
+def __builtin_amdgcn_image_store_1d_v4f16_i32 : AMDGPUBuiltin<"void(_ExtVector<4, _Float16>, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts,16-bit-insts">;
 def __builtin_amdgcn_image_store_1darray_v4f32_i32 : AMDGPUBuiltin<"void(_ExtVector<4, float>, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
-def __builtin_amdgcn_image_store_1darray_v4f16_i32 : AMDGPUBuiltin<"void(_ExtVector<4, _Float16>, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
+def __builtin_amdgcn_image_store_1darray_v4f16_i32 : AMDGPUBuiltin<"void(_ExtVector<4, _Float16>, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts,16-bit-insts">;
 def __builtin_amdgcn_image_store_2d_f32_i32 : AMDGPUBuiltin<"void(float, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
 def __builtin_amdgcn_image_store_2d_v4f32_i32 : AMDGPUBuiltin<"void(_ExtVector<4, float>, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
-def __builtin_amdgcn_image_store_2d_v4f16_i32 : AMDGPUBuiltin<"void(_ExtVector<4, _Float16>, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
+def __builtin_amdgcn_image_store_2d_v4f16_i32 : AMDGPUBuiltin<"void(_ExtVector<4, _Float16>, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts,16-bit-insts">;
 def __builtin_amdgcn_image_store_2darray_f32_i32 : AMDGPUBuiltin<"void(float, int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
 def __builtin_amdgcn_image_store_2darray_v4f32_i32 : AMDGPUBuiltin<"void(_ExtVector<4, float>, int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
-def __builtin_amdgcn_image_store_2darray_v4f16_i32 : AMDGPUBuiltin<"void(_ExtVector<4, _Float16>, int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
+def __builtin_amdgcn_image_store_2darray_v4f16_i32 : AMDGPUBuiltin<"void(_ExtVector<4, _Float16>, int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts,16-bit-insts">;
 def __builtin_amdgcn_image_store_3d_v4f32_i32 : AMDGPUBuiltin<"void(_ExtVector<4, float>, int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
-def __builtin_amdgcn_image_store_3d_v4f16_i32 : AMDGPUBuiltin<"void(_ExtVector<4, _Float16>, int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
+def __builtin_amdgcn_image_store_3d_v4f16_i32 : AMDGPUBuiltin<"void(_ExtVector<4, _Float16>, int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts,16-bit-insts">;
 def __builtin_amdgcn_image_store_cube_v4f32_i32 : AMDGPUBuiltin<"void(_ExtVector<4, float>, int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
-def __builtin_amdgcn_image_store_cube_v4f16_i32 : AMDGPUBuiltin<"void(_ExtVector<4, _Float16>, int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
+def __builtin_amdgcn_image_store_cube_v4f16_i32 : AMDGPUBuiltin<"void(_ExtVector<4, _Float16>, int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts,16-bit-insts">;
 def __builtin_amdgcn_image_store_mip_1d_v4f32_i32 : AMDGPUBuiltin<"void(_ExtVector<4, float>, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
-def __builtin_amdgcn_image_store_mip_1d_v4f16_i32 : AMDGPUBuiltin<"void(_ExtVector<4, _Float16>, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
+def __builtin_amdgcn_image_store_mip_1d_v4f16_i32 : AMDGPUBuiltin<"void(_ExtVector<4, _Float16>, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts,16-bit-insts">;
 def __builtin_amdgcn_image_store_mip_1darray_v4f32_i32 : AMDGPUBuiltin<"void(_ExtVector<4, float>, int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
-def __builtin_amdgcn_image_store_mip_1darray_v4f16_i32 : AMDGPUBuiltin<"void(_ExtVector<4, _Float16>, int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
+def __builtin_amdgcn_image_store_mip_1darray_v4f16_i32 : AMDGPUBuiltin<"void(_ExtVector<4, _Float16>, int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts,16-bit-insts">;
 def __builtin_amdgcn_image_store_mip_2d_f32_i32 : AMDGPUBuiltin<"void(float, int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
 def __builtin_amdgcn_image_store_mip_2d_v4f32_i32 : AMDGPUBuiltin<"void(_ExtVector<4, float>, int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
-def __builtin_amdgcn_image_store_mip_2d_v4f16_i32 : AMDGPUBuiltin<"void(_ExtVector<4, _Float16>, int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
+def __builtin_amdgcn_image_store_mip_2d_v4f16_i32 : AMDGPUBuiltin<"void(_ExtVector<4, _Float16>, int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts,16-bit-insts">;
 def __builtin_amdgcn_image_store_mip_2darray_f32_i32 : AMDGPUBuiltin<"void(float, int, int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
 def __builtin_amdgcn_image_store_mip_2darray_v4f32_i32 : AMDGPUBuiltin<"void(_ExtVector<4, float>, int, int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
-def __builtin_amdgcn_image_store_mip_2darray_v4f16_i32 : AMDGPUBuiltin<"void(_ExtVector<4, _Float16>, int, int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
+def __builtin_amdgcn_image_store_mip_2darray_v4f16_i32 : AMDGPUBuiltin<"void(_ExtVector<4, _Float16>, int, int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts,16-bit-insts">;
 def __builtin_amdgcn_image_store_mip_3d_v4f32_i32 : AMDGPUBuiltin<"void(_ExtVector<4, float>, int, int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
-def __builtin_amdgcn_image_store_mip_3d_v4f16_i32 : AMDGPUBuiltin<"void(_ExtVector<4, _Float16>, int, int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
+def __builtin_amdgcn_image_store_mip_3d_v4f16_i32 : AMDGPUBuiltin<"void(_ExtVector<4, _Float16>, int, int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts,16-bit-insts">;
 def __builtin_amdgcn_image_store_mip_cube_v4f32_i32 : AMDGPUBuiltin<"void(_ExtVector<4, float>, int, int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
-def __builtin_amdgcn_image_store_mip_cube_v4f16_i32 : AMDGPUBuiltin<"void(_ExtVector<4, _Float16>, int, int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts">;
+def __builtin_amdgcn_image_store_mip_cube_v4f16_i32 : AMDGPUBuiltin<"void(_ExtVector<4, _Float16>, int, int, int, int, int, __amdgpu_texture_t, int, int)", [Const], "image-insts,16-bit-insts">;
 def __builtin_amdgcn_image_sample_1d_v4f32_f32 : AMDGPUBuiltin<"_ExtVector<4, float>(int, float, __amdgpu_texture_t, _ExtVector<4, int>, bool, int, int)", [Const], "image-insts">;
 def __builtin_amdgcn_image_sample_1d_v4f16_f32 : AMDGPUBuiltin<"_ExtVector<4, _Float16>(int, float, __amdgpu_texture_t, _ExtVector<4, int>, bool, int, int)", [Const], "image-insts">;
 def __builtin_amdgcn_image_sample_1darray_v4f32_f32 : AMDGPUBuiltin<"_ExtVector<4, float>(int, float, float, __amdgpu_texture_t, _ExtVector<4, int>, bool, int, int)", [Const], "image-insts">;
diff --git a/clang/test/Sema/builtins-amdgcn-d16-image-16bit-error.c b/clang/test/Sema/builtins-amdgcn-d16-image-16bit-error.c
new file mode 100644
index 0000000000000..9f4f18e572177
--- /dev/null
+++ b/clang/test/Sema/builtins-amdgcn-d16-image-16bit-error.c
@@ -0,0 +1,59 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx700 -verify -fsyntax-only %s
+
+// Verify that half typed image intrinsics require 16-bit-insts as well as
+// image-insts.
+
+typedef _Float16 half;
+typedef half half4 __attribute__((ext_vector_type(4)));
+
+void test(half4 v, __amdgpu_texture_t tex) {
+  v = __builtin_amdgcn_image_load_1d_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+      0, 0, tex, 0, 0);
+  v = __builtin_amdgcn_image_load_1darray_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+      0, 0, 0, tex, 0, 0);
+  v = __builtin_amdgcn_image_load_2d_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+      0, 0, 0, tex, 0, 0);
+  v = __builtin_amdgcn_image_load_2darray_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+      0, 0, 0, 0, tex, 0, 0);
+  v = __builtin_amdgcn_image_load_3d_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+      0, 0, 0, 0, tex, 0, 0);
+  v = __builtin_amdgcn_image_load_cube_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+      0, 0, 0, 0, tex, 0, 0);
+  v = __builtin_amdgcn_image_load_mip_1d_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+      0, 0, 0, tex, 0, 0);
+  v = __builtin_amdgcn_image_load_mip_1darray_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+      0, 0, 0, 0, tex, 0, 0);
+  v = __builtin_amdgcn_image_load_mip_2d_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+      0, 0, 0, 0, tex, 0, 0);
+  v = __builtin_amdgcn_image_load_mip_2darray_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+      0, 0, 0, 0, 0, tex, 0, 0);
+  v = __builtin_amdgcn_image_load_mip_3d_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+      0, 0, 0, 0, 0, tex, 0, 0);
+  v = __builtin_amdgcn_image_load_mip_cube_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+      0, 0, 0, 0, 0, tex, 0, 0);
+  __builtin_amdgcn_image_store_1d_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+      v, 0, 0, tex, 0, 0);
+  __builtin_amdgcn_image_store_1darray_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+      v, 0, 0, 0, tex, 0, 0);
+  __builtin_amdgcn_image_store_2d_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+      v, 0, 0, 0, tex, 0, 0);
+  __builtin_amdgcn_image_store_2darray_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+      v, 0, 0, 0, 0, tex, 0, 0);
+  __builtin_amdgcn_image_store_3d_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+      v, 0, 0, 0, 0, tex, 0, 0);
+  __builtin_amdgcn_image_store_cube_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+      v, 0, 0, 0, 0, tex, 0, 0);
+  __builtin_amdgcn_image_store_mip_1d_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+      v, 0, 0, 0, tex, 0, 0);
+  __builtin_amdgcn_image_store_mip_1darray_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+      v, 0, 0, 0, 0, tex, 0, 0);
+  __builtin_amdgcn_image_store_mip_2d_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+      v, 0, 0, 0, 0, tex, 0, 0);
+  __builtin_amdgcn_image_store_mip_2darray_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+      v, 0, 0, 0, 0, 0, tex, 0, 0);
+  __builtin_amdgcn_image_store_mip_3d_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+      v, 0, 0, 0, 0, 0, tex, 0, 0);
+  __builtin_amdgcn_image_store_mip_cube_v4f16_i32( // expected-error {{needs target feature image-insts,16-bit-insts}}
+      v, 0, 0, 0, 0, 0, tex, 0, 0);
+}

arsenm · 2026-06-23T16:00:51Z

Requires ROCm#3033 to land first

ranapratap55 · 2026-06-27T14:05:30Z

Please add similar changes to sample/gather4 cases.

arsenm mentioned this pull request Jun 23, 2026

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

Open

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

arsenm requested review from ranapratap55 and shiltian June 23, 2026 16:00

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

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

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

clang/AMDGPU: Require 16-bit-insts for half typed image builtins#205368
arsenm wants to merge 1 commit into
users/arsenm/clang/amdgpu-fix-accepting-d16-format-builtins-gfx6-gfx7from
users/arsenm/clang/amdgpu-fix-accepting-d16-image-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!

arsenm commented Jun 23, 2026

Uh oh!

ranapratap55 commented Jun 27, 2026 •

edited

Loading

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 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!

arsenm commented Jun 23, 2026

Uh oh!

ranapratap55 commented Jun 27, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

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

edited

Loading

ranapratap55 commented Jun 27, 2026 •

edited

Loading