From feaf31184bc20619448c87f3c38a1bbc3ec21719 Mon Sep 17 00:00:00 2001 From: Aaditya Date: Tue, 30 Sep 2025 11:37:42 +0530 Subject: [PATCH] [AMDGPU] Add builtins for wave reduction intrinsics --- clang/include/clang/Basic/BuiltinsAMDGPU.def | 4 ++++ clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp | 4 ++++ 2 files changed, 8 insertions(+) diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def index fda16e42d2c6b..ebc0ac35f42d9 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.def +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def @@ -402,6 +402,10 @@ BUILTIN(__builtin_amdgcn_wave_reduce_max_u64, "WUiWUiZi", "nc") BUILTIN(__builtin_amdgcn_wave_reduce_and_b64, "WiWiZi", "nc") BUILTIN(__builtin_amdgcn_wave_reduce_or_b64, "WiWiZi", "nc") BUILTIN(__builtin_amdgcn_wave_reduce_xor_b64, "WiWiZi", "nc") +BUILTIN(__builtin_amdgcn_wave_reduce_add_f32, "ffZi", "nc") +BUILTIN(__builtin_amdgcn_wave_reduce_sub_f32, "ffZi", "nc") +BUILTIN(__builtin_amdgcn_wave_reduce_min_f32, "ffZi", "nc") +BUILTIN(__builtin_amdgcn_wave_reduce_max_f32, "ffZi", "nc") //===----------------------------------------------------------------------===// // R600-NI only builtins. diff --git a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp index 07cf08c54985a..a242a73e4a822 100644 --- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp +++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp @@ -301,18 +301,22 @@ static Intrinsic::ID getIntrinsicIDforWaveReduction(unsigned BuiltinID) { llvm_unreachable("Unknown BuiltinID for wave reduction"); case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_add_u32: case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_add_u64: + case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_add_f32: return Intrinsic::amdgcn_wave_reduce_add; case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_sub_u32: case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_sub_u64: + case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_sub_f32: return Intrinsic::amdgcn_wave_reduce_sub; case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_min_i32: case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_min_i64: + case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_min_f32: return Intrinsic::amdgcn_wave_reduce_min; case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_min_u32: case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_min_u64: return Intrinsic::amdgcn_wave_reduce_umin; case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_max_i32: case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_max_i64: + case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_max_f32: return Intrinsic::amdgcn_wave_reduce_max; case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_max_u32: case clang::AMDGPU::BI__builtin_amdgcn_wave_reduce_max_u64: