[llvm-branch-commits] [clang] clang/AMDGPU: Add half typed image gather4 builtin (PR #213613)
Matt Arsenault via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Mon Aug 3 01:31:27 PDT 2026
https://github.com/arsenm updated https://github.com/llvm/llvm-project/pull/213613
>From 5d90a4da93e5d096412fdb661c50ad28b3b5e6a5 Mon Sep 17 00:00:00 2001
From: Matt Arsenault <Matthew.Arsenault at amd.com>
Date: Mon, 3 Aug 2026 08:24:55 +0200
Subject: [PATCH] clang/AMDGPU: Add half typed image gather4 builtin
Follow along with the precedent of using an f32 suffix
for the coordinate type. We probably should have had one
builtin that detected the coordinate type.
Co-Authored-By: Claude (Opus 4.8) <noreply at anthropic.com>
---
clang/include/clang/Basic/BuiltinsAMDGPU.td | 1 +
clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 1 +
clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp | 1 +
clang/lib/Sema/SemaAMDGPU.cpp | 8 +++--
.../builtins-amdgcn-extended-image.hip | 8 +++++
clang/test/CodeGen/builtins-extended-image.c | 30 +++++++++++++++++++
.../builtins-amdgcn-d16-image-16bit-error.c | 5 ++++
.../SemaOpenCL/builtins-extended-image-err.cl | 5 ++++
...iltins-extended-image-param-gfx1100-err.cl | 15 ++++++++++
9 files changed, 72 insertions(+), 2 deletions(-)
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.td b/clang/include/clang/Basic/BuiltinsAMDGPU.td
index cc8e80b482c44..9bea517d7b888 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.td
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.td
@@ -1404,3 +1404,4 @@ def __builtin_amdgcn_image_sample_d_2darray_v4f16_f32 : AMDGPUBuiltin<"_ExtVecto
def __builtin_amdgcn_image_sample_d_3d_v4f32_f32 : AMDGPUBuiltin<"_ExtVector<4, float>(int, float, float, float, float, float, float, float, float, float, __amdgpu_texture_t, _ExtVector<4, int>, bool, int, int)", [Const], "extended-image-insts">;
def __builtin_amdgcn_image_sample_d_3d_v4f16_f32 : AMDGPUBuiltin<"_ExtVector<4, _Float16>(int, float, float, float, float, float, float, float, float, float, __amdgpu_texture_t, _ExtVector<4, int>, bool, int, int)", [Const], "extended-image-insts,16-bit-insts">;
def __builtin_amdgcn_image_gather4_lz_2d_v4f32_f32 : AMDGPUBuiltin<"_ExtVector<4, float>(int, float, float, __amdgpu_texture_t, _ExtVector<4, int>, bool, int, int)", [Const], "extended-image-insts">;
+def __builtin_amdgcn_image_gather4_lz_2d_v4f16_f32 : AMDGPUBuiltin<"_ExtVector<4, _Float16>(int, float, float, __amdgpu_texture_t, _ExtVector<4, int>, bool, int, int)", [Const], "extended-image-insts,16-bit-insts">;
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index e2a1c2b94dec4..c3ebfda38e728 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -730,6 +730,7 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
return emitAMDGCNImageOverloadedReturnType(
*this, expr, "amdgcn.image.sample.d.2darray", false);
case AMDGPU::BI__builtin_amdgcn_image_gather4_lz_2d_v4f32_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_gather4_lz_2d_v4f16_f32:
return emitAMDGCNImageOverloadedReturnType(
*this, expr, "amdgcn.image.gather4.lz.2d", false);
case AMDGPU::BI__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4:
diff --git a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
index 29199f1726c1a..c2aa02453f885 100644
--- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
+++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
@@ -1374,6 +1374,7 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
return emitAMDGCNImageOverloadedReturnType(
*this, E, Intrinsic::amdgcn_image_sample_d_2darray, false);
case clang::AMDGPU::BI__builtin_amdgcn_image_gather4_lz_2d_v4f32_f32:
+ case clang::AMDGPU::BI__builtin_amdgcn_image_gather4_lz_2d_v4f16_f32:
return emitAMDGCNImageOverloadedReturnType(
*this, E, Intrinsic::amdgcn_image_gather4_lz_2d, false);
case AMDGPU::BI__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4:
diff --git a/clang/lib/Sema/SemaAMDGPU.cpp b/clang/lib/Sema/SemaAMDGPU.cpp
index 48230fa262d5c..a0602024e3122 100644
--- a/clang/lib/Sema/SemaAMDGPU.cpp
+++ b/clang/lib/Sema/SemaAMDGPU.cpp
@@ -267,7 +267,8 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,
case AMDGPU::BI__builtin_amdgcn_image_sample_d_2darray_v4f16_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_d_3d_v4f32_f32:
case AMDGPU::BI__builtin_amdgcn_image_sample_d_3d_v4f16_f32:
- case AMDGPU::BI__builtin_amdgcn_image_gather4_lz_2d_v4f32_f32: {
+ case AMDGPU::BI__builtin_amdgcn_image_gather4_lz_2d_v4f32_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_gather4_lz_2d_v4f16_f32: {
StringRef FeatureList(
getASTContext().BuiltinInfo.getRequiredFeatures(BuiltinID));
if (!Builtin::evaluateRequiredTargetFeatures(FeatureList,
@@ -307,7 +308,10 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,
// For gather, only one bit can be set indicating which exact component to
// return.
bool ExtraGatherChecks =
- BuiltinID == AMDGPU::BI__builtin_amdgcn_image_gather4_lz_2d_v4f32_f32 &&
+ (BuiltinID ==
+ AMDGPU::BI__builtin_amdgcn_image_gather4_lz_2d_v4f32_f32 ||
+ BuiltinID ==
+ AMDGPU::BI__builtin_amdgcn_image_gather4_lz_2d_v4f16_f32) &&
SemaRef.BuiltinConstantArgPower2(TheCall, 0);
return ExtraGatherChecks ||
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-extended-image.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-extended-image.hip
index 32c51594e380e..7153ea82bc78e 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-extended-image.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-extended-image.hip
@@ -50,6 +50,14 @@ __device__ float4 test_gather4_lz_2d_v4f32_a(float s, float t, __amdgpu_texture_
return __builtin_amdgcn_image_gather4_lz_2d_v4f32_f32(8, s, t, tex, samp, 0, 120, 110);
}
+// CIR-LABEL: @_Z26test_gather4_lz_2d_v4f16_r
+// CIR: cir.call_llvm_intrinsic "amdgcn.image.gather4.lz.2d" {{.*}} : (!s32i, !cir.float, !cir.float, !cir.vector<8 x !s32i>, !cir.vector<4 x !s32i>, !cir.bool, !s32i, !s32i) -> !cir.vector<4 x !cir.f16>
+// LLVM: define{{.*}} <4 x half> @_Z26test_gather4_lz_2d_v4f16_rffu18__amdgpu_texture_tDv4_i(
+// LLVM: call {{.*}}<4 x half> @llvm.amdgcn.image.gather4.lz.2d.v4f16.f32.v8i32.v4i32(i32 1, float {{.*}}, float {{.*}}, <8 x i32> {{.*}}, <4 x i32> {{.*}}, i1 {{.*}}, i32 {{.*}}, i32 {{.*}})
+__device__ half4 test_gather4_lz_2d_v4f16_r(float s, float t, __amdgpu_texture_t tex, int4 samp) {
+ return __builtin_amdgcn_image_gather4_lz_2d_v4f16_f32(1, s, t, tex, samp, 0, 120, 110);
+}
+
// CIR-LABEL: @_Z23test_sample_lz_1d_v4f32
// CIR: cir.call_llvm_intrinsic "amdgcn.image.sample.lz.1d" {{.*}} : (!s32i, !cir.float, !cir.vector<8 x !s32i>, !cir.vector<4 x !s32i>, !cir.bool, !s32i, !s32i) -> !cir.vector<4 x !cir.float>
// LLVM: define{{.*}} <4 x float> @_Z23test_sample_lz_1d_v4f32fu18__amdgpu_texture_tDv4_i(
diff --git a/clang/test/CodeGen/builtins-extended-image.c b/clang/test/CodeGen/builtins-extended-image.c
index 42c2bfd360174..4c1e85edb4cb6 100644
--- a/clang/test/CodeGen/builtins-extended-image.c
+++ b/clang/test/CodeGen/builtins-extended-image.c
@@ -125,6 +125,36 @@ float4 test_amdgcn_image_gather4_lz_2d_v4f32_f32_a(float4 v4f32, float f32, int
return __builtin_amdgcn_image_gather4_lz_2d_v4f32_f32(8, f32, f32, tex, vec4i32, 0, 120, 110);
}
+// CHECK-LABEL: define dso_local <4 x half> @test_amdgcn_image_gather4_lz_2d_v4f16_f32_r(
+// CHECK-SAME: <4 x half> noundef [[V4F16:%.*]], float noundef [[F32:%.*]], i32 noundef [[I32:%.*]], <8 x i32> [[TEX:%.*]], <4 x i32> noundef [[VEC4I32:%.*]]) #[[ATTR0]] {
+// CHECK-NEXT: [[ENTRY:.*:]]
+// CHECK-NEXT: [[V4F16_ADDR:%.*]] = alloca <4 x half>, align 8, addrspace(5)
+// CHECK-NEXT: [[F32_ADDR:%.*]] = alloca float, align 4, addrspace(5)
+// CHECK-NEXT: [[I32_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: [[TEX_ADDR:%.*]] = alloca <8 x i32>, align 32, addrspace(5)
+// CHECK-NEXT: [[VEC4I32_ADDR:%.*]] = alloca <4 x i32>, align 16, addrspace(5)
+// CHECK-NEXT: [[V4F16_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[V4F16_ADDR]] to ptr
+// CHECK-NEXT: [[F32_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[F32_ADDR]] to ptr
+// CHECK-NEXT: [[I32_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[I32_ADDR]] to ptr
+// CHECK-NEXT: [[TEX_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[TEX_ADDR]] to ptr
+// CHECK-NEXT: [[VEC4I32_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[VEC4I32_ADDR]] to ptr
+// CHECK-NEXT: store <4 x half> [[V4F16]], ptr [[V4F16_ADDR_ASCAST]], align 8
+// CHECK-NEXT: store float [[F32]], ptr [[F32_ADDR_ASCAST]], align 4
+// CHECK-NEXT: store i32 [[I32]], ptr [[I32_ADDR_ASCAST]], align 4
+// CHECK-NEXT: store <8 x i32> [[TEX]], ptr [[TEX_ADDR_ASCAST]], align 32
+// CHECK-NEXT: store <4 x i32> [[VEC4I32]], ptr [[VEC4I32_ADDR_ASCAST]], align 16
+// CHECK-NEXT: [[TMP0:%.*]] = load float, ptr [[F32_ADDR_ASCAST]], align 4
+// CHECK-NEXT: [[TMP1:%.*]] = load float, ptr [[F32_ADDR_ASCAST]], align 4
+// CHECK-NEXT: [[TMP2:%.*]] = load <8 x i32>, ptr [[TEX_ADDR_ASCAST]], align 32
+// CHECK-NEXT: [[TMP3:%.*]] = load <4 x i32>, ptr [[VEC4I32_ADDR_ASCAST]], align 16
+// CHECK-NEXT: [[TMP4:%.*]] = call <4 x half> @llvm.amdgcn.image.gather4.lz.2d.v4f16.f32.v8i32.v4i32(i32 1, float [[TMP0]], float [[TMP1]], <8 x i32> [[TMP2]], <4 x i32> [[TMP3]], i1 false, i32 120, i32 110)
+// CHECK-NEXT: ret <4 x half> [[TMP4]]
+//
+half4 test_amdgcn_image_gather4_lz_2d_v4f16_f32_r(half4 v4f16, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) {
+
+ return __builtin_amdgcn_image_gather4_lz_2d_v4f16_f32(1, f32, f32, tex, vec4i32, 0, 120, 110);
+}
+
// CHECK-LABEL: define dso_local <4 x float> @test_amdgcn_image_sample_lz_1d_v4f32_f32(
// CHECK-SAME: <4 x float> noundef [[V4F32:%.*]], float noundef [[F32:%.*]], i32 noundef [[I32:%.*]], <8 x i32> [[TEX:%.*]], <4 x i32> noundef [[VEC4I32:%.*]]) #[[ATTR0]] {
// CHECK-NEXT: [[ENTRY:.*:]]
diff --git a/clang/test/Sema/builtins-amdgcn-d16-image-16bit-error.c b/clang/test/Sema/builtins-amdgcn-d16-image-16bit-error.c
index 4f7fcb4fc9cea..c7680b907f37e 100644
--- a/clang/test/Sema/builtins-amdgcn-d16-image-16bit-error.c
+++ b/clang/test/Sema/builtins-amdgcn-d16-image-16bit-error.c
@@ -107,3 +107,8 @@ void test_sample(half4 v, float f32, __amdgpu_texture_t tex, int4 vec4i32) {
v = __builtin_amdgcn_image_sample_d_3d_v4f16_f32( // expected-error {{needs target feature extended-image-insts,16-bit-insts}}
15, f32, f32, f32, f32, f32, f32, f32, f32, f32, tex, vec4i32, 0, 0, 0);
}
+
+void test_gather4(half4 v, float f32, __amdgpu_texture_t tex, int4 vec4i32) {
+ v = __builtin_amdgcn_image_gather4_lz_2d_v4f16_f32( // expected-error {{needs target feature extended-image-insts,16-bit-insts}}
+ 1, f32, f32, tex, vec4i32, 0, 0, 0);
+}
diff --git a/clang/test/SemaOpenCL/builtins-extended-image-err.cl b/clang/test/SemaOpenCL/builtins-extended-image-err.cl
index 24cf0308cfcf0..f295d03d93fa1 100644
--- a/clang/test/SemaOpenCL/builtins-extended-image-err.cl
+++ b/clang/test/SemaOpenCL/builtins-extended-image-err.cl
@@ -30,6 +30,11 @@ float4 test_amdgcn_image_gather4_lz_2d_v4f32_f32_a(float4 v4f32, float f32, int
return __builtin_amdgcn_image_gather4_lz_2d_v4f32_f32(1, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_gather4_lz_2d_v4f32_f32_a' needs target feature extended-image-insts}}
}
+half4 test_amdgcn_image_gather4_lz_2d_v4f16_f32_r(half4 v4f16, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) {
+
+ return __builtin_amdgcn_image_gather4_lz_2d_v4f16_f32(1, f32, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_gather4_lz_2d_v4f16_f32_r' needs target feature extended-image-insts}}
+}
+
float4 test_amdgcn_image_sample_lz_1d_v4f32_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) {
return __builtin_amdgcn_image_sample_lz_1d_v4f32_f32(15, f32, tex, vec4i32, 0, 101, 121); //GFX94-error{{'test_amdgcn_image_sample_lz_1d_v4f32_f32' needs target feature extended-image-insts}}
diff --git a/clang/test/SemaOpenCL/builtins-extended-image-param-gfx1100-err.cl b/clang/test/SemaOpenCL/builtins-extended-image-param-gfx1100-err.cl
index 0610d3facf2f6..0f85f99f736d6 100644
--- a/clang/test/SemaOpenCL/builtins-extended-image-param-gfx1100-err.cl
+++ b/clang/test/SemaOpenCL/builtins-extended-image-param-gfx1100-err.cl
@@ -37,6 +37,21 @@ float4 test_amdgcn_image_gather4_lz_2d_v4f32_f32_dmask_range(float4 v4f32, float
return __builtin_amdgcn_image_gather4_lz_2d_v4f32_f32(16, f32, f32, tex, vec4i32, 0, 120, 110); //expected-error{{argument value 16 is outside the valid range [0, 15]}}
}
+half4 test_amdgcn_image_gather4_lz_2d_v4f16_f32_r(half4 v4f16, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) {
+
+ return __builtin_amdgcn_image_gather4_lz_2d_v4f16_f32(1, f32, f32, tex, vec4i32, 0, f32, i32); //expected-error{{argument to '__builtin_amdgcn_image_gather4_lz_2d_v4f16_f32' must be a constant integer}}
+}
+
+half4 test_amdgcn_image_gather4_lz_2d_v4f16_f32_dmask_power_of_2(half4 v4f16, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) {
+
+ return __builtin_amdgcn_image_gather4_lz_2d_v4f16_f32(3, f32, f32, tex, vec4i32, 0, 120, 110); //expected-error{{argument should be a power of 2}}
+}
+
+half4 test_amdgcn_image_gather4_lz_2d_v4f16_f32_dmask_range(half4 v4f16, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) {
+
+ return __builtin_amdgcn_image_gather4_lz_2d_v4f16_f32(16, f32, f32, tex, vec4i32, 0, 120, 110); //expected-error{{argument value 16 is outside the valid range [0, 15]}}
+}
+
float4 test_amdgcn_image_sample_lz_1d_v4f32_f32(float4 v4f32, float f32, int i32, __amdgpu_texture_t tex, int4 vec4i32) {
return __builtin_amdgcn_image_sample_lz_1d_v4f32_f32(i32, f32, tex, vec4i32, 0, f32, i32); //expected-error{{argument to '__builtin_amdgcn_image_sample_lz_1d_v4f32_f32' must be a constant integer}}
More information about the llvm-branch-commits
mailing list