[clang] [llvm] [AMDGPU] Add builtins for wave reduction intrinsics (PR #127013)
Matt Arsenault via cfe-commits
cfe-commits at lists.llvm.org
Tue Feb 18 22:32:59 PST 2025
================
@@ -346,6 +346,24 @@ BUILTIN(__builtin_amdgcn_endpgm, "v", "nr")
BUILTIN(__builtin_amdgcn_get_fpenv, "WUi", "n")
BUILTIN(__builtin_amdgcn_set_fpenv, "vWUi", "n")
+//===----------------------------------------------------------------------===//
+
+// Wave Reduction builtins.
+
+//===----------------------------------------------------------------------===//
+
+BUILTIN(__builtin_amdgcn_wave_reduce_add_i32, "ii", "nc")
+BUILTIN(__builtin_amdgcn_wave_reduce_uadd_i32, "ii", "nc")
----------------
arsenm wrote:
uadd_i32 -> add_u32 would be the more consistent naming convention. You're also using signed integer types for all of these, should use unsigned for the unsigned operations
https://github.com/llvm/llvm-project/pull/127013
More information about the cfe-commits
mailing list