__shfl_down_sync and the mask argument

This is the warp reduction from NVIDIA's own post on warp-level primitives, with the membership mask computed exactly the way that post says to compute it.

unsigned mask = __ballot_sync(FULL_MASK, threadIdx.x < NUM_ELEMENTS);
if (threadIdx.x < NUM_ELEMENTS) {
    val = input[threadIdx.x];
    for (int offset = 16; offset > 0; offset /= 2)
        val += __shfl_down_sync(mask, val, offset);
}

(https://developer.nvidia.com/blog/using-cuda-warp-level-primitives/ , checked 2026-08-30.)

Every rule about the mask holds here: each calling lane has its bit set, each non-calling lane has its bit clear, everyone passes the same value. Set NUM_ELEMENTS to 20 and the answer is still undefined, because at the first step lane 4 asks for lane 20, which is not in that mask and is not running that line. The guide carries an invalid example of the same shape.

The mask names the threads that must turn up. It does not make a thread you did not name readable. Those are two different questions, and the loop above only answers the first.

By the end you can write a warp sum and a warp max that use no shared memory and no barrier, and say of any shuffle whether its mask is enough.

Lanes reading each other's registers

A warp shuffle is one instruction that moves a value from one lane of a warp to another. It needs no address or allocation, and it does not send the value through memory. The guide states the rule: "Warp shuffle functions exchange a value between non-exited threads within a warp without the use of shared memory."

There are four of them, and the difference is only how each lane works out which lane to read from:

T __shfl_sync     (unsigned mask, T value, int      srcLane,  int width=warpSize);
T __shfl_up_sync  (unsigned mask, T value, unsigned delta,    int width=warpSize);
T __shfl_down_sync(unsigned mask, T value, unsigned delta,    int width=warpSize);
T __shfl_xor_sync (unsigned mask, T value, int      laneMask, int width=warpSize);

__shfl_sync reads a lane you name, which is a broadcast when every lane names the same one. __shfl_down_sync reads lane L + delta and __shfl_up_sync reads L - delta, and neither wraps: a lane with no partner keeps the value it had. __shfl_xor_sync reads lane L xor laneMask, which gives every lane a partner at every step.

That last difference decides the shape of your reduction. Five __shfl_down_sync steps at deltas 16, 8, 4, 2 and 1 leave the total in lane 0 and nowhere else, so if the rest of the warp wants it you pay a sixth instruction to broadcast it back. Five __shfl_xor_sync steps leave the answer in all 32.

Neither is faster. Pick the one whose answer ends up where you want it.

The registers make this cheap. The value never leaves the register file, so it needs no store, load, or address. The exchange happens inside one instruction, so you need no __syncthreads(), __syncwarp(), or shared array.

A block-wide reduction still needs a barrier where warps combine, which is day 24. A reduction within one warp does not.

The mask names who shows up, not who you can read

The intuition you arrive with is that the mask is a required argument you fill in, the way you fill in a stream argument, and that 0xffffffff means "the usual thing". It is the set of lanes that have to reach this instruction, and the hardware waits for them.

The guide is direct about it. Each calling thread must have its bit set, and each non-calling thread must have its bit clear. All non-exited threads named in the mask must execute the intrinsic with the same mask.

Break those rules and the behaviour is "invalid, such as kernel hang, or undefined". A common failure is passing 0xffffffff from inside a branch that half the warp took. Warp divergence usually costs throughput, but here it can stop the kernel from returning.

Now the part that catches people who did compute a mask. A shuffle reads a lane, and there is a second rule for the lane it reads: "Threads may only read data from another thread that is actively participating in the intrinsics. If the target thread is inactive, the retrieved value is undefined." The guide's own invalid example is a correct mask over five lanes:

if (laneId <= 4) {
    // undefined behavior: destination lanes 5, 6 are not active for lanes 3, 4
    result = __shfl_down_sync(0b11111, value, 2);
}

So 0xffffffff is right only when all 32 lanes are there. A narrower mask is not the usual fix when some lanes have no data. Keep every lane inside the collective and give the unused lanes the operation's identity value.

That is day 14's barrier rule applied to a warp operation. Guard the loads and stores, never the collective. Since compute capability 7.0, independent thread scheduling has given each lane its own program counter, so these operations use _sync and a mask instead of assuming convergence.

One warp, one row, three ways

Full program in code/day23-shuffles/shuffles.cu. One warp reduces one row of a matrix. A row holds 20 live values inside a 32-column stride, so the group that owns data is smaller than the warp and is not a power of two, which is the only case where any of this matters.

Three rules hold it together, and the first is what earns the right to pass 0xffffffff.

Every lane stays inside every collective. The lanes with no data carry the identity: 0 for the sum, a large negative float for the max, false for __any_sync, true for __all_sync. That is what earns the right to pass 0xffffffff.

The comparison kernel is the fair one. reduceRowsShared does the same reduction through a shared tile with __syncwarp(), not __syncthreads(), because a per-warp reduction never needed a block barrier. Comparing against a block barrier would add work that the shuffle kernel does not need.

The claim is measured by the driver, not asserted. cudaFuncGetAttributes reports static shared memory per block for each kernel. The page's central claim is that number, and the program prints it rather than checking it, because a gate on the claim under test makes the test unfalsifiable.

The sum, in five steps and a broadcast:

    // Five steps, each halving the number of lanes that still hold a partial
    // sum. Lane L reads lane L + offset, so after the last step lane 0 holds
    // the total and no other lane does.
    float sum = value;
    for (unsigned int offset = kWarpSize / 2; offset > 0; offset /= 2) {
        sum += __shfl_down_sync(kFullMask, sum, offset);
    }

    // Every lane needs the total to normalise its own value, and only lane 0
    // has it, so one broadcast copies lane 0's register into all 32.
    const float rowSum = __shfl_sync(kFullMask, sum, 0);

The max, in five steps and no broadcast:

    // The same five steps as a butterfly. Lane L swaps with lane L xor
    // laneMask, so every lane ends up holding the max and no broadcast is
    // needed. Same step count as the ladder above, different lane owns the
    // answer at the end, and that is the whole difference between the two.
    float best = live ? value : kMaxIdentity;
    for (int laneMask = kWarpSize / 2; laneMask > 0; laneMask /= 2) {
        const float other = __shfl_xor_sync(kFullMask, best, laneMask);
        best = (other > best) ? other : best;
    }

The votes are the other half of the instruction set, and they collapse a per-lane predicate into one warp-wide answer:

    // Three questions about the whole warp, one instruction each. The lanes
    // with no data vote the identity of the question they are in, which is
    // how all 32 can stay named in the mask without changing an answer.
    const unsigned int liveMask = __ballot_sync(kFullMask, live);
    const int liveCount = __popc(liveMask);
    const int anyHigh = __any_sync(kFullMask, live && value > kHighWater);
    const int allPositive = __all_sync(kFullMask, !live || value > 0.0f);

__ballot_sync returns a bit per lane, __any_sync and __all_sync return the OR and the AND of the same predicate. A ballot is how you compute a mask for a later collective, and __popc on it is how you count, both in one instruction.

The third kernel makes the faulty case runnable:

    // Every lane of the warp reaches the ballot, so kFullMask is right here.
    // liveMask comes back as 0x000fffff, lanes 0 to 19, which is a legal mask
    // and is what the canonical example tells you to pass.
    const unsigned int liveMask = __ballot_sync(kFullMask, live);
    if (live) {
        float sum = in[row * kStride + lane];
        for (unsigned int offset = kWarpSize / 2; offset > 0; offset /= 2) {
            // Legal mask, undefined read. At offset 16 lane 4 asks for lane
            // 20, which liveMask does not name and which is not running this
            // line, so what comes back is whatever that lane's register
            // happens to hold.
            sum += __shfl_down_sync(liveMask, sum, offset);
        }
        if (lane == 0) {
            sums[row] = sum;
        }
    }

Note. That kernel ships and the other broken one does not, and the difference is the failure mode. This one satisfies every mask rule and breaks only the source-lane rule, so it returns a bad number. Passing 0xffffffff from inside a branch breaks a mask rule, which the guide documents as able to hang, and a build that hangs teaches nothing a paragraph cannot.

Results

Measured. Tesla T4, driver 595.84, CUDA 12.6 (V12.6.85), built with nvcc -O3 -arch=sm_75. Captured 2026-08-30 on the project's verification node; full transcript in the page's evidence file.

GPU: Tesla T4 (compute capability 7.5)
warpSize from the driver: 32, compiled against: 32
rows = 8205, 20 live columns in a 32 column stride, 256 threads per block, 1026 blocks
8208 warps launched, 3 of them own no row

Kernel attributes, from cudaFuncGetAttributes
kernel                   shared/block  registers   local/thread
reduceRowsPadded                    0         23              0
reduceRowsShared                 2048         15              0
reduceRowsBallotDown                0         16              0

Correctness against the CPU reference, 8205 rows
kernel                  sums wrong  maxes wrong
reduceRowsPadded                 0            0
reduceRowsShared                 0            0
reduceRowsBallotDown             0            -
  every row came out right, which is a fact about this driver and not about the code

reduceRowsPadded against reduceRowsShared: 0 of 8205 rows differ bit for bit
normalised elements outside a 1e-05 relative tolerance: 0 of 262560
__popc(__ballot_sync(...)) disagrees with 20 on 0 rows
__any_sync and __all_sync disagree with the reference on 0 rows

every gated kernel matched the CPU reference

A shuffle reduction uses no shared memory at all. reduceRowsShared reserves 2048 bytes per block; reduceRowsPadded and reduceRowsBallotDown reserve zero and reach the same answers on all 8205 rows. On a card with 48 KiB per block, shared memory is the resource that limits how many blocks fit, so a reduction that does not need it leaves that budget for something else.

The register counts are worth reading beside it: 23, 15 and 16. The shuffle versions are not free, they move the cost from shared memory into registers, which is usually the trade you want.

The correctness line says something careful. Every row came out right, which is a fact about this driver and this input, not a proof that the masks are correct. A shuffle with a wrong mask can produce right answers for a long time and then fail when a warp is partially active. The masks here are correct by construction, because every collective is reached by all 32 lanes.

Run it yourself

Compiler Explorer, which is free and executes on a Tesla T4. The program allocates two 1 MiB buffers and launches three kernels once each, so it finishes well inside the 20 second run cap that rules most of this course's programs out. Target sm_75 or lower and pin the compiler to a specific nvcc entry rather than trunk.

The file is about 690 lines, past this site's 60-line embed budget, so there is no box on this page yet. Paste it into https://godbolt.org/ (checked 2026-08-30) and the whole program runs. A shorter embed with one correct kernel and the faulty kernel over a few rows still needs to be built and pinned to a specific nvcc id.

Colab's free T4 and any card you own work too. The build line is in the repo's README:

nvcc -std=c++17 -O3 -arch=sm_75 -o shuffles shuffles.cu

Exercise

Add a sixth output to reduceRowsPadded: argmax[row], the lowest lane index that holds the row's maximum. Use only warp primitives, no shared memory and no second pass over the row.

Time: 20 to 30 minutes. Submit: your shuffles.cu and the one sentence in the solution's last paragraph, in your own words.

Check: the harness runs the same 8,205 rows and compares your argmax column against a CPU reference that breaks ties by the lowest column. It reports the smallest failing row, not the first one it meets, and prints your value, the reference value and that row's twenty values, so you can see which tie you lost. Ties are what separate a working answer from one that happens to pass, and this input has none, so the harness runs a second case built to contain them.

const unsigned int holders = __ballot_sync(kFullMask, live && value == best);
const int argmax = __ffs(static_cast<int>(holders)) - 1;

Pitfalls

Your kernel hangs and the last thing you changed was a shuffle. You passed 0xffffffff from inside a branch only some lanes reach, so the hardware waits for lanes that never arrive. Compute the mask with __ballot_sync before the branch, or restructure so every lane arrives. Day 22 covers the branches that split a warp.

You used __activemask() as the mask. It reports which lanes happen to be converged at that instant, not which lanes your algorithm needs, and the guide says it "cannot be used to determine which warp lanes execute a given branch". NVIDIA's own advice is one line: "Don't just use __activemask() as the mask value." Day 21 measured what it does report.

Your mask is correct and your answer is still wrong. A sequence of __shfl_down_sync calls reads lane L + delta, and every lane it reads has to be in the mask too. Over a group that is not the whole warp it is not, and the value is undefined. Keep all 32 lanes in and pad with the identity.

You dropped the _sync suffix because an older example did. The legacy primitives without a mask are deprecated from CUDA 9.0, and the current guide documents only the _sync forms. They relied on lanes staying in lockstep, which stopped being guaranteed at compute capability 7.0.

You wrapped an old primitive in __syncwarp() and called it fixed. After the threads leave __syncwarp() they are free to diverge again, so the shuffle underneath still runs with lanes missing. Convergence is guaranteed inside the primitive and nowhere else.

You reduced a whole block with shuffles alone. A shuffle reaches 32 lanes and there is no shuffle between warps. Reduce inside each warp, write one value per warp, then reduce those. Day 24 compares the block reduction versions and shows where the barrier returns.

Go deeper

• CUDA Programming Guide 5.4.6 "Warp Functions", and 5.4.6.6 "Warp __sync Intrinsic Constraints" for the mask rules: https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/cpp-language-extensions.html (checked 2026-08-30)
• "Using CUDA Warp-Level Primitives": https://developer.nvidia.com/blog/using-cuda-warp-level-primitives/ (checked 2026-08-30)
• cuda-samples, cpp/2_Concepts_and_Techniques/shfl_scan: https://github.com/NVIDIA/cuda-samples/tree/master/cpp/2_Concepts_and_Techniques/shfl_scan (checked 2026-08-30)
• Programming Massively Parallel Processors, 4th edition, chapter 10, on reduction: https://shop.elsevier.com/books/programming-massively-parallel-processors/hwu/978-0-323-91231-0

Next

Day 24 turns these five steps into a block-wide reduction and starts timing versions against each other. Day 25 ends the comparison with a shuffle version.

Past one warp the guarantees change again: cooperative groups on day 28 names the group you want rather than making you track masks by hand, and CUB, which the course reaches on day 39, ships WarpReduce so you never write this loop in production.