llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-backend-amdgpu 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 _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
