[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