src/utils/sync.cuh

0001 /** TRACCC library, part of the ACTS project (R&D line)
0002  *
0003  * (c) 2021-2022 CERN for the benefit of the ACTS project
0004  *
0005  * Mozilla Public License Version 2.0
0006  */
0007
0008 #pragma once
0009
0010 namespace traccc::cuda {
0011 /**
0012  * @brief Calculate the index of the current thread in the set of threads in
0013  * the warp for which a predicate evaluates to true.
0014  *
0015  * @param predicate The predicate to evaluate.
0016  * @param mask The thread mask to sync.
0017  *
0018  * @returns The tuple (T, I), where T is the total number of threads in the
0019  * warp for which the provided predicate was true, and where I is a unique,
0020  * consecutive index in the set of threads for which this was true.
0021  *
0022  * @example If a predicate P evaluates as `[T, F, F, T]` in a 4-lane warp, the
0023  * call to `warp_indexed_ballot_sync(P)` will evaluate to the values
0024  * `[(2, 0), (2, ?), (2, ?), (2, 1)]` where `?` indicates an undefined value.
0025  *
0026  * @warning The value of T is always well-defined, the value if I is only
0027  * well-defined if the given predicate was true for the given thread.
0028  *
0029  * @note As with all CUDA synchronization barriers, exited threads are treated
0030  * as having implicitly reached the barrier to avoid deadlock situations.
0031  *
0032  * @note This function forces thread synchronization.
0033  */
0034 __device__ __forceinline__ std::pair<uint32_t, uint32_t>
0035 warp_indexed_ballot_sync(bool predicate, uint32_t mask = 0xFFFFFFFFu) {
0036   uint32_t vote = __ballot_sync(mask, predicate);
0037
0038   /*
0039    * The total number of threads which return true is simply the population
0040    * count of the voting result, which is to say the number of true bits in
0041    * its binary expansion.
0042    */
0043   uint32_t tot = __popc(vote);
0044
0045   /*
0046    * The index is the population count of the vote mask, but only for bits
0047    * with an index lower than the current thread index. Thus, we shift a bit
0048    * mask over the voting result, nulling any bits that are _higher_ than
0049    * the thread index!
0050    */
0051   uint32_t idx = __popc(vote & ~(0xFFFFFFFFu << (threadIdx.x % warpSize)));
0052
0053   return {tot, idx};
0054 }
0055 }  // namespace traccc::cuda