Predicates, or truth values, can be represented compactly, since each
occupies only 1 bit. CUDA includes several instructions that make
reducing predicates efficient. Within a warp,
__ballot_sync() and __popc() reduce a
per-thread condition to a count:
unsigned __ballot_sync( unsigned mask, int predicate );
int __popc( unsigned v );__ballot_sync() evaluates predicate for
every participating thread and returns a 32-bit word whose bits give the
condition for each lane. Because it broadcasts that word to every lane,
it is effectively a reduction across the warp; __popc()
then counts the set bits. At block scope,
__syncthreads_count() does both steps at once:
int __syncthreads_count( int predicate );It waits until all warps in the block have arrived, then returns to
every thread the number of threads for which predicate was
true. Because the 1-bit predicates turn into 5-bit (warp) or up to
10-bit (block) counts immediately, these intrinsics mainly reduce the
amount of shared memory needed for the lowest level of the reduction –
but they greatly increase the number of elements a single thread block
can consider.