Warp-Level Collective Decisions with `__any_sync` and `__all_sync`

The Principle of Warp-Level Voting

At its core, a vote function is a collective operation where each participating thread in a warp contributes a boolean value, known as a predicate. The function then aggregates these predicates and broadcasts a single, unified result back to all participating threads. This is analogous to a committee vote: each member votes 'yes' or 'no', and the outcome (e.g., "was the vote unanimous?") is announced to everyone at once.

The two primary vote functions are:

• int __any_sync(unsigned mask, int predicate);: This function returns true (a non-zero integer) if the predicate is true for at least one of the threads specified in the mask. Otherwise, it returns false (zero).
• int __all_sync(unsigned mask, int predicate);: This function returns true if the predicate is true for all of the threads specified in the mask. Otherwise, it returns false.

Crucially, like the shuffle primitives, these functions are synchronous. They ensure all threads designated in the mask have reached the call and evaluated their predicate before the collective result is computed and returned. This explicit synchronization is the key to their safety and correctness.

The following presentation slide deck, which we've referenced before, provides a concise summary of the principles and categories of these modern, synchronous built-ins.

A Practical Example: Conditional Algorithm Selection

Vote functions are particularly powerful for optimizing code where a group of threads must collectively decide on a strategy. Imagine a scenario where you have multiple versions of an algorithm—for instance, a highly optimized memory copy for 16-byte aligned data and a more general, slower version for unaligned data.

You want the entire warp to use the fast path if and only if every thread in the warp has properly aligned data. Using a simple if statement on each thread could lead to divergence: some threads take the fast path, others the slow path, negating the performance benefit. __all_sync provides an elegant solution.

The presentation slides you just reviewed contain an excellent example of this.

Let's break down the logic from the slides:

__device__ void memorycopy(char *tptr, char *sptr, size_t size) {
    // 1. Get a mask of all currently active threads in the warp.
    unsigned mask = __activemask();

    // 2. Check if all active threads have 16-byte aligned pointers.
    // The predicate 'is_all_aligned' is assumed to be a device function
    // that returns true if both pointers are aligned.
    if (__all_sync(mask, is_all_aligned(tptr, sptr, 16))) {
        // If the vote passes (unanimous 'true'), all threads
        // execute the specialized 16-byte copy. No divergence.
        return memcpy_aligned_16(tptr, sptr, size);
    }

    // 3. If the first vote failed, try the next best option.
    if (__all_sync(mask, is_all_aligned(tptr, sptr, 8))) {
        return memcpy_aligned_8(tptr, sptr, size);
    }

    // ... and so on for other alignments, or a default case.
}

This pattern is extremely efficient. The vote is a fast, hardware-supported operation. It prevents divergence by ensuring all participating threads agree on a single execution path, which aligns perfectly with your interest in low-level performance systems.

__ballot_sync: A Detailed Vote Tally

While __any_sync and __all_sync give a simple "yes/no" summary, __ballot_sync provides a much more detailed result.

• unsigned int __ballot_sync(unsigned mask, int predicate);: This function also takes a predicate from each participating thread. However, instead of returning a single boolean, it returns a 32-bit unsigned int. Each bit in this integer corresponds to a lane in the warp (bit i corresponds to lane i). The bit is set to 1 if the thread in that lane had a true predicate and was part of the input mask. Otherwise, the bit is 0.

__ballot_sync is like a roll-call vote. It doesn't just tell you if someone voted 'yes'; it gives you a bitmask identifying everyone who voted 'yes'.

This is incredibly useful for dynamically creating participation masks for other warp-level operations. We saw this in the previous lesson's resources when preparing for a shuffle-based reduction. It's worth revisiting this now with a deeper understanding.

The example from the article is a perfect illustration of combining primitives:

// Assume NUM_ELEMENTS could be less than the block size.
// We want to create a mask for only the threads that will do work.
unsigned mask = __ballot_sync(FULL_MASK, threadIdx.x < NUM_ELEMENTS);

if (threadIdx.x < NUM_ELEMENTS) {
    val = input[threadIdx.x];
    // This reduction now uses the precise mask of participating threads.
    // If a thread tries to shuffle from a lane where the mask bit is 0,
    // the behavior is well-defined (it gets a 0).
    for (int offset = 16; offset > 0; offset /= 2) {
        val += __shfl_down_sync(mask, val, offset);
    }
    …
}

Here, __ballot_sync acts as a "scout." It runs ahead, polls the threads based on a condition, and returns a map of valid participants. This map is then fed into the subsequent shuffle operation, ensuring it operates correctly even at the boundary of the dataset.

Conclusion

In this lesson, you've added another critical set of tools to your low-level CUDA toolkit. Where shuffle primitives manage data exchange, vote primitives manage collective decision-making. By understanding and using these functions, you can write code that is not only faster but also more robust and often more expressive.

Key Takeaways:

• Warp vote functions (__any_sync, __all_sync) allow threads in a warp to make fast, collective boolean decisions, which is ideal for avoiding control flow divergence.
• __ballot_sync provides a more detailed outcome, returning a bitmask that identifies exactly which participating threads satisfied a given predicate.
• The mask argument is fundamental to the correct and safe use of these primitives, making the set of participating threads explicit and removing dangerous reliance on implicit compiler behavior.
• These primitives are composable: __ballot_sync is often used to generate a precise mask for subsequent shuffle or vote operations, enabling complex and efficient intra-warp coordination.

We have now established a solid understanding of intra-warp communication and coordination. In the next module, we will begin applying these and other techniques to implement and optimize fundamental parallel algorithms, or "primitives," like reduction and scan, which form the building blocks of almost all complex GPU workloads.