[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