[libc-commits] [libc] [libc] Add utility functions for warp-level scan and reduction (PR #84866)

Joseph Huber via libc-commits libc-commits at lists.llvm.org
Tue Mar 12 05:10:07 PDT 2024


================
@@ -31,6 +31,25 @@ LIBC_INLINE bool is_first_lane(uint64_t lane_mask) {
   return gpu::get_lane_id() == get_first_lane_id(lane_mask);
 }
 
+/// Gets the sum of all lanes inside the warp or wavefront.
----------------
jhuber6 wrote:

I definitely see the appeal of having these as intrinsic functions. Specifically because right now I need to use `__builtin_amdgcn_wavefrontsize` which prevents these loops from being fully unrolled. However, I think it will be easy to drop those intrinsic functions in later.

https://github.com/llvm/llvm-project/pull/84866


More information about the libc-commits mailing list