[llvm-branch-commits] [clang] clang/AMDGPU: Add half typed image gather4 builtin (PR #213613)
via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Mon Aug 3 00:26:18 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clang
Author: Matt Arsenault (arsenm)
<details>
<summary>Changes</summary>
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@<!-- -->anthropic.com>
---
Full diff: https://github.com/llvm/llvm-project/pull/213613.diff
8 Files Affected:
- (modified) clang/include/clang/Basic/BuiltinsAMDGPU.td (+1)
- (modified) clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp (+1)
- (modified) clang/lib/Sema/SemaAMDGPU.cpp (+6-2)
- (modified) clang/test/CIR/CodeGenHIP/builtins-amdgcn-extended-image.hip (+8)
- (modified) clang/test/CodeGen/builtins-extended-image.c (+30)
- (modified) clang/test/Sema/builtins-amdgcn-d16-image-16bit-error.c (+5)
- (modified) clang/test/SemaOpenCL/builtins-extended-image-err.cl (+5)
- (modified) clang/test/SemaOpenCL/builtins-extended-image-param-gfx1100-err.cl (+15)
``````````diff
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/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}}
``````````
</details>
https://github.com/llvm/llvm-project/pull/213613
More information about the llvm-branch-commits
mailing list