Skip to content
Markdown

Warp-level CUDA primitives: shuffle, vote, match, and syncwarp

Scope: the _sync warp-level intrinsics, shuffle (__shfl_sync, __shfl_up_sync, __shfl_down_sync, __shfl_xor_sync), vote (__all_sync, __any_sync, __ballot_sync), match (__match_any_sync, __match_all_sync), __activemask(), and __syncwarp(), that move data and predicates between the 32 lanes of a warp through the register file instead of shared memory. Covers the mandatory mask argument, the undefined-behavior contract it enforces since Volta's independent thread scheduling, and why the deprecated non-_sync forms are unsafe on current hardware.

The CUDA C++ blocks below are reference templates: they need nvcc and a GPU to compile and run, and are labeled as such. The numpy block in Validated (numpy, runnable) is self-contained, executed, and asserted in this page.

What it is

A warp is 32 threads issued in lockstep under SIMT (see GPU Execution Model: SM, Warp, SIMT). Warp-level primitives let the 32 lanes of a warp exchange register values and combine per-lane predicates directly, without staging through shared memory and a __syncthreads() barrier. Three families cover this:

  • Shuffle (__shfl_sync, __shfl_up_sync, __shfl_down_sync, __shfl_xor_sync): each lane reads another named lane's register value. __shfl_sync broadcasts a fixed source lane; __shfl_up_sync/__shfl_down_sync shift by a lane offset (used for scans and reductions); __shfl_xor_sync swaps with lane ^ laneMask (a butterfly exchange, used for reductions that leave the result in every lane).
  • Vote (__all_sync, __any_sync, __ballot_sync): each lane contributes a boolean predicate; the warp reduces it to a single answer. __all_sync/__any_sync return a scalar (AND/OR across the named lanes); __ballot_sync returns a 32-bit mask with bit N set when lane N's predicate was true, useful for counting or for building a mask to pass to another intrinsic.
  • Match (__match_any_sync, __match_all_sync): each lane contributes a value; the warp groups lanes by equality. __match_any_sync returns, per lane, the bitmask of lanes sharing its value; __match_all_sync returns the full mask (and sets *pred true) only if every named lane holds the identical value.
  • __activemask() returns the set of lanes currently executing the calling instruction, as a 32-bit mask.
  • __syncwarp(mask) synchronizes and orders memory operations among the threads named in mask; every named thread must reach the call.

Every one of these is a _sync intrinsic: it takes an explicit unsigned mask naming its participants. This replaced an older, implicit-warp-sync family (__shfl, __ballot, __any, __all, no _sync suffix, no mask argument) that assumed all 32 lanes of a warp always moved together. That assumption broke with Volta's independent thread scheduling (ITS): each thread got its own program counter, so a warp's lanes are no longer guaranteed to reconverge at the same point a pre-Volta compiler assumed. The non-_sync forms are deprecated and their implicit full-warp behavior is unsafe on ITS hardware; current CUDA documents only the _sync forms.1

Why use it

  • No shared memory, no barrier. Data moves lane-to-lane through the register file in one instruction, versus a shared-memory write, a __syncthreads(), and a shared-memory read. This removes both the barrier latency and the shared-memory traffic for small, warp-scoped exchanges (reductions, scans, broadcasts, deduplication).
  • The mask makes divergence explicit. Requiring every intrinsic call to name its participants forces the programmer (or the compiler, via __activemask()) to state exactly which lanes are expected to reach the call, which is the information ITS hardware needs to run divergent code correctly.
  • Cheap building blocks for common warp idioms. Warp-level reductions (shuffle), leader election and active-lane counting (vote), and warp-aggregated atomics / key deduplication (match) are all one or two intrinsics instead of hand-rolled shared-memory code.

When to use it (and when not)

  • Use warp shuffle for a reduction or scan whose scope is one warp (or width-wide sub-groups of a warp), for a broadcast of one lane's value to the rest of the warp, or for a warp-level transpose, instead of routing through shared memory.
  • Use warp vote to test a condition across the warp cheaply (e.g. "does any lane in this warp still have work") or to build a mask (__ballot_sync) for a subsequent shuffle or atomic.
  • Use warp match to deduplicate identical keys/addresses within a warp before an atomic (warp-aggregated atomics), cutting atomic contention when many lanes update the same location.
  • Use __syncwarp() after code that leaves a warp's lanes at different points (e.g. following divergent branches or before/after volatile/shared-memory accesses that need warp-wide visibility) on architectures with ITS, where reconvergence is not automatic.
  • Do not reach for these primitives above the warp: they cannot communicate across warps or blocks (use shared memory + __syncthreads(), or global memory + a grid-wide sync, for that). See Shared Memory Tiling for the cross-warp equivalent.
  • Do not use the non-_sync (__shfl, __ballot, __any, __all) forms in new code: they lack a mask, are deprecated, and their implicit-full-warp assumption is unsafe on Volta-and-later ITS hardware.
  • Do not hardcode 0xFFFFFFFF out of habit. Only use the full mask when the call site is genuinely reached by all 32 lanes unconditionally; inside any conditional or loop with a data-dependent exit, compute the real mask.

Architecture

flowchart TB
  subgraph WARP["One warp = 32 lanes, one instruction issued per cycle (SIMT)"]
    direction LR
    L0["Lane 0 reg"]
    L1["Lane 1 reg"]
    L2["Lane 2 reg"]
    DOTS["..."]
    L31["Lane 31 reg"]
  end
  MASK["mask: unsigned 32-bit\nbit N set = lane N is a named participant"]
  MASK -->|"every named lane MUST execute\nthe SAME intrinsic with the SAME mask"| WARP

  SHFL["__shfl_sync / _up / _down / _xor\nlane reads another named lane's register"]
  VOTE["__ballot_sync / __all_sync / __any_sync\nreduce per-lane predicates to one bitmask/bool"]
  MATCH["__match_any_sync / __match_all_sync\ngroup named lanes with equal value"]
  SYNCW["__syncwarp(mask)\norders memory ops, reconverges named lanes"]

  WARP --> SHFL
  WARP --> VOTE
  WARP --> MATCH
  WARP --> SYNCW

  SHFL -->|"register-to-register"| L2
  VOTE -->|"one bitmask/bool back to every named lane"| L1
  MATCH -->|"per-lane group bitmask"| L0

  SMEM["Alternative: shared memory\nwrite -> __syncthreads() -> read"]
  WARP -.->|"what these intrinsics let you skip\nfor warp-scoped exchanges"| SMEM

The mask is not a hint: it is the contract. Hardware and the compiler assume every lane named in mask reaches that exact call with that exact mask value; anything that violates it (a stale mask computed before a divergent branch, or a mask naming a lane that took the other branch) is undefined behavior, not a slower-but-correct fallback.

How it works

Reference: warp shuffle-down reduction (full warp)

The canonical warp-sum reduction. mask here is 0xFFFFFFFF because this call site is reached unconditionally by the full warp; the result lands in lane 0 (and the other lanes whose local index plus the running offset stays in range).

// Reference template: requires nvcc + a CUDA-capable GPU. Not executed in this KB.
__device__ float warp_reduce_sum(float val) {
    for (int offset = 16; offset > 0; offset >>= 1) {
        val += __shfl_down_sync(0xFFFFFFFF, val, offset);
    }
    return val;  // full sum in lane 0 when the whole warp reached this call
}

Reference: computing the mask under divergence

The unsafe pattern is calling a _sync intrinsic inside a diverged branch with a mask computed before the divergence, or with a hardcoded 0xFFFFFFFF. The safe pattern computes the mask at the point of the call, from the predicate that caused the divergence, with __ballot_sync seeded from the warp's currently active lanes:

// Reference template: requires nvcc + a CUDA-capable GPU. Not executed in this KB.
__device__ void process_if_valid(int* data, int idx, int n) {
    bool valid = idx < n;                              // this predicate causes divergence
    unsigned active = __activemask();                  // lanes reaching this line right now
    unsigned valid_mask = __ballot_sync(active, valid); // mask of the lanes that are valid

    if (valid) {
        // Every lane inside this branch is named in valid_mask, and every one of
        // them reaches this call, so this __shfl_down_sync is well-defined.
        int v = data[idx];
        int partner_sum = v + __shfl_down_sync(valid_mask, v, 1, 32);
        data[idx] = partner_sum;
    }
    // WRONG (do not do this): __shfl_down_sync(0xFFFFFFFF, v, 1) inside the `if`
    // above. Lanes that failed `valid` never reach the call, so 0xFFFFFFFF names
    // lanes that are not there: undefined behavior, observed as either a hang
    // (Volta+ can wait on a lane that never arrives) or a corrupted read.
}

Reference: vote, match, and syncwarp

// Reference template: requires nvcc + a CUDA-capable GPU. Not executed in this KB.
__device__ void warp_aggregated_histogram(int* bins, int key) {
    unsigned mask = __activemask();                 // lanes live at this call
    unsigned peers = __match_any_sync(mask, key);    // lanes in *this* warp sharing `key`
    int leader = __ffs(peers) - 1;                   // lowest-numbered peer lane, 1 atomic per group
    int count = __popc(peers);
    if ((threadIdx.x % 32) == leader) {
        atomicAdd(&bins[key], count);                // one atomic instead of one per lane
    }
    __syncwarp(mask);                                // reconverge the named lanes before continuing
}

__device__ bool warp_any_out_of_bounds(int idx, int n) {
    unsigned mask = __activemask();
    return __any_sync(mask, idx >= n) != 0;
}

Validated (numpy, runnable)

The block below models all four intrinsic families, mask semantics, and partial-warp behavior as plain numpy, with no CUDA dependency. It was executed with python3 before being pasted here; every assert passed (terminal output in the report below). The adversarial case at the end reproduces the exact failure the "WRONG" comment above describes: assuming a full 0xFFFFFFFF mask when only some lanes hold valid data silently corrupts a reduction instead of raising an error.

"""
Software model of warp-level CUDA intrinsics (shuffle, vote, match) plus the
mask/partial-warp semantics that make the _sync intrinsics correct on Volta+
independent thread scheduling. Pure numpy, no CUDA required.

Every lane is one element of a 32-length numpy array. A "mask" is a 32-bit
Python int, bit N set means lane N is a named participant, exactly like the
`unsigned mask` argument of __shfl_sync / __ballot_sync / __match_any_sync /
__syncwarp. The hardware contract (CUDA C++ Programming Guide, Warp Shuffle
Functions / Warp Vote Functions / Warp Match Functions): every thread named
in mask must execute the SAME intrinsic call with the SAME mask value, or the
result is undefined. This model does not (and cannot) enforce that at the
Python level -- it is a caller discipline, exactly like on real hardware --
but it does model the two ways real hardware punishes violating it:
(1) reading a shuffle source lane that is not in mask is undefined (modeled
    as NaN, a visibly wrong value, rather than silently 0), and
(2) a reduction that runs over a mask wider than the truly active lanes
    silently folds in stale/garbage data from inactive lanes.
"""
import numpy as np

WARP_SIZE = 32
FULL_MASK = 0xFFFFFFFF


def mask_from_lanes(active_lanes):
    m = 0
    for lane in active_lanes:
        m |= (1 << lane)
    return m


def lane_active(mask, lane):
    return bool((mask >> lane) & 1)


def shfl_down_sync(mask, values, delta, width=WARP_SIZE):
    """Model of T __shfl_down_sync(unsigned mask, T var, unsigned delta, int width=warpSize).
    Lane L reads from lane L+delta within its width-sized segment. Reading from a
    lane outside mask is undefined in hardware; here it surfaces as NaN so a bad
    read is visible instead of silently defaulting to 0."""
    out = np.full(WARP_SIZE, np.nan)
    for lane in range(WARP_SIZE):
        if not lane_active(mask, lane):
            continue
        seg_base = (lane // width) * width
        src = lane + delta
        if src >= seg_base + width:
            out[lane] = values[lane]
        elif not lane_active(mask, src):
            out[lane] = np.nan
        else:
            out[lane] = values[src]
    return out


def shfl_up_sync(mask, values, delta, width=WARP_SIZE):
    """Model of T __shfl_up_sync(unsigned mask, T var, unsigned delta, int width=warpSize)."""
    out = np.full(WARP_SIZE, np.nan)
    for lane in range(WARP_SIZE):
        if not lane_active(mask, lane):
            continue
        seg_base = (lane // width) * width
        src = lane - delta
        if src < seg_base:
            out[lane] = values[lane]
        elif not lane_active(mask, src):
            out[lane] = np.nan
        else:
            out[lane] = values[src]
    return out


def shfl_xor_sync(mask, values, lane_mask_xor, width=WARP_SIZE):
    """Model of T __shfl_xor_sync(unsigned mask, T var, int laneMask, int width=warpSize):
    butterfly exchange, lane L swaps with lane (L ^ laneMask) inside its segment."""
    out = np.full(WARP_SIZE, np.nan)
    for lane in range(WARP_SIZE):
        if not lane_active(mask, lane):
            continue
        seg_base = (lane // width) * width
        local = lane - seg_base
        src = seg_base + (local ^ lane_mask_xor)
        if not lane_active(mask, src):
            out[lane] = np.nan
        else:
            out[lane] = values[src]
    return out


def shfl_sync(mask, values, src_lane, width=WARP_SIZE):
    """Model of T __shfl_sync(unsigned mask, T var, int srcLane, int width=warpSize):
    every lane in its segment reads the same srcLane (broadcast within segment)."""
    out = np.full(WARP_SIZE, np.nan)
    for lane in range(WARP_SIZE):
        if not lane_active(mask, lane):
            continue
        seg_base = (lane // width) * width
        src = seg_base + (src_lane % width)
        if not lane_active(mask, src):
            out[lane] = np.nan
        else:
            out[lane] = values[src]
    return out


def ballot_sync(mask, predicate):
    """Model of unsigned __ballot_sync(unsigned mask, int predicate): bit N set
    iff lane N is in mask AND predicate[N] is nonzero."""
    result = 0
    for lane in range(WARP_SIZE):
        if lane_active(mask, lane) and predicate[lane]:
            result |= (1 << lane)
    return result


def all_sync(mask, predicate):
    """Model of int __all_sync(unsigned mask, int predicate)."""
    for lane in range(WARP_SIZE):
        if lane_active(mask, lane) and not predicate[lane]:
            return 0
    return 1


def any_sync(mask, predicate):
    """Model of int __any_sync(unsigned mask, int predicate)."""
    for lane in range(WARP_SIZE):
        if lane_active(mask, lane) and predicate[lane]:
            return 1
    return 0


def match_any_sync(mask, values):
    """Model of unsigned __match_any_sync(unsigned mask, T value): per lane in
    mask, bitmask of lanes in mask sharing the same value (self included)."""
    out = np.zeros(WARP_SIZE, dtype=np.uint32)
    for lane in range(WARP_SIZE):
        if not lane_active(mask, lane):
            continue
        m = 0
        for other in range(WARP_SIZE):
            if lane_active(mask, other) and values[other] == values[lane]:
                m |= (1 << other)
        out[lane] = m
    return out


def match_all_sync(mask, values):
    """Model of unsigned __match_all_sync(unsigned mask, T value, int *pred):
    returns (result, pred) arrays. pred[L]=1 and result[L]=mask iff every
    active lane in mask holds the identical value; else pred[L]=0."""
    active = [l for l in range(WARP_SIZE) if lane_active(mask, l)]
    uniform = len({values[l] for l in active}) == 1
    result = np.zeros(WARP_SIZE, dtype=np.uint32)
    pred = np.zeros(WARP_SIZE, dtype=np.int32)
    for lane in active:
        pred[lane] = 1 if uniform else 0
        result[lane] = mask if uniform else 0
    return result, pred


def warp_reduce_sum_full(values, mask=FULL_MASK):
    """Canonical full-warp shuffle-down sum reduction:
    for (offset = 16; offset > 0; offset /= 2) val += __shfl_down_sync(mask, val, offset);
    Correct when every lane named in mask is truly active (the common case:
    a full 32-wide warp). Lane 0 ends up holding the full sum when mask == FULL_MASK."""
    v = values.astype(np.float64).copy()
    offset = 16
    while offset > 0:
        shuffled = shfl_down_sync(mask, v, offset)
        v = v + np.where(np.isnan(shuffled), 0.0, shuffled)
        offset //= 2
    return v


def warp_reduce_sum_partial_correct(values, active_lanes):
    """Correct partial-warp reduction: build a mask restricted to the actually
    active (contiguous, 0-based) lanes, and bound each shuffle step by the
    active count so a lane never folds in a partner outside the group. This
    is the documented safe pattern for tail iterations with fewer than 32
    live lanes (e.g. the last chunk of a grid-stride loop)."""
    n = len(active_lanes)
    mask = mask_from_lanes(active_lanes)
    v = values.astype(np.float64).copy()
    offset = 1
    while offset < n:
        shfl_down_sync(mask, v, offset)  # models the real __shfl_down_sync call site
        for i, lane in enumerate(active_lanes):
            if i + offset < n:
                partner_lane = active_lanes[i + offset]
                v[lane] = v[lane] + v[partner_lane]
        offset *= 2
    return v, mask


def warp_reduce_sum_buggy_full_mask(values, active_lanes):
    """Adversarial / buggy version: the reduction was written assuming a full,
    fully-populated 32-lane warp, so it hardcodes FULL_MASK and runs the
    unconditional textbook loop (offset 16 down to 1) with no notion that
    only `active_lanes` hold real data. On real hardware the SIMT warp still
    executes uniformly (no divergent branch here, so no hang), but lanes
    outside the true active set hold stale/garbage register contents (e.g.
    an over-read past the tail of a grid-stride loop). Since every lane is
    named in FULL_MASK, __shfl_down_sync happily folds that garbage into the
    sum: the mask lied about who has valid data, and the result is silently
    wrong instead of undefined-behavior-flagged."""
    active_set = set(active_lanes)
    v = np.full(WARP_SIZE, np.nan)
    for lane in active_lanes:
        v[lane] = values[lane]
    rng = np.random.default_rng(42)
    for lane in range(WARP_SIZE):
        if lane not in active_set:
            v[lane] = rng.uniform(100.0, 200.0)  # stale/garbage, never meant to be read
    reduced = warp_reduce_sum_full(v, FULL_MASK)
    return reduced[active_lanes[0]]


# ---------------------------------------------------------------------------
# Tests
# ---------------------------------------------------------------------------

def slow_reduce(values, active_lanes):
    return sum(values[l] for l in active_lanes)


# 1. Full-warp shuffle-down sum reduction, mask = FULL_MASK, all 32 lanes active
values32 = np.arange(1, 33, dtype=np.float64)  # 1..32
result = warp_reduce_sum_full(values32, FULL_MASK)
expected_total = values32.sum()
assert np.isclose(result[0], expected_total), (result[0], expected_total)
print(f"full-warp reduce: lane0={result[0]}, expected={expected_total}")

# 2. shfl_xor_sync butterfly reduction should also land the full sum in every lane
xor_mask = FULL_MASK
xor_offset = 16
vv = values32.copy()
while xor_offset > 0:
    shuffled = shfl_xor_sync(xor_mask, vv, xor_offset)
    vv = vv + shuffled
    xor_offset //= 2
assert np.allclose(vv, expected_total), vv
print("butterfly (xor) reduce: all lanes hold", vv[0], "expected", expected_total)

# 3. shfl_sync broadcast: every active lane reads lane 5's value
bcast_mask = mask_from_lanes(range(32))
bcast = shfl_sync(bcast_mask, values32, 5)
assert np.all(bcast == values32[5]), bcast
print("shfl_sync broadcast of lane 5:", bcast[0], bcast[31])

# 4. shfl_up_sync: lane L reads lane L-1 (inclusive prefix shift), lane 0 keeps its own
up = shfl_up_sync(bcast_mask, values32, 1)
assert up[0] == values32[0]
assert np.allclose(up[1:], values32[:-1])
print("shfl_up_sync delta=1 ok")

# 5. Vote functions: predicate true for even-valued lanes
predicate = (values32.astype(int) % 2 == 0).astype(int)
full_mask_all = mask_from_lanes(range(32))
assert all_sync(full_mask_all, predicate) == 0          # not all lanes are even
assert any_sync(full_mask_all, predicate) == 1           # some lanes are even
ballot = ballot_sync(full_mask_all, predicate)
expected_ballot = mask_from_lanes([l for l in range(32) if predicate[l]])
assert ballot == expected_ballot, (bin(ballot), bin(expected_ballot))
print("vote functions: all_sync=0, any_sync=1, ballot matches expected bitmask")

all_true_predicate = np.ones(32, dtype=int)
assert all_sync(full_mask_all, all_true_predicate) == 1
print("all_sync=1 when every active lane's predicate is true")

# 6. Vote functions restricted to a partial mask (only lanes 0..12 participate,
# a non-power-of-two group of 13 lanes)
partial_lanes = list(range(13))
partial_mask = mask_from_lanes(partial_lanes)
pred13 = (values32.astype(int) % 2 == 0).astype(int)
ballot13 = ballot_sync(partial_mask, pred13)
assert ballot13 & ~partial_mask == 0, "ballot leaked bits outside the mask"
expected13 = mask_from_lanes([l for l in partial_lanes if pred13[l]])
assert ballot13 == expected13, (bin(ballot13), bin(expected13))
print("partial-mask (13 lanes) ballot_sync correctly scoped, no bit leakage")

# 7. Match functions: three groups of matching values across 10 active lanes
match_values = np.array([7, 7, 3, 3, 3, 9, 7, 3, 9, 1] + [0] * 22, dtype=np.int64)
match_lanes = list(range(10))
match_mask = mask_from_lanes(match_lanes)
match_result = match_any_sync(match_mask, match_values)
assert match_result[0] == mask_from_lanes([0, 1, 6])     # lanes with value 7
assert match_result[2] == mask_from_lanes([2, 3, 4, 7])  # lanes with value 3
assert match_result[5] == mask_from_lanes([5, 8])        # lanes with value 9
assert match_result[9] == mask_from_lanes([9])           # lane 9 (value 1) matches only itself
print("match_any_sync groups lanes by value correctly")

uniform_values = np.array([42] * 8 + [0] * 24, dtype=np.int64)
uniform_mask = mask_from_lanes(range(8))
res_all, pred_all = match_all_sync(uniform_mask, uniform_values)
assert np.all(pred_all[:8] == 1)
assert np.all(res_all[:8] == uniform_mask)
print("match_all_sync: uniform group -> pred=1, result=mask")

nonuniform_values = np.array([42] * 7 + [43] + [0] * 24, dtype=np.int64)
res_nu, pred_nu = match_all_sync(uniform_mask, nonuniform_values)
assert np.all(pred_nu[:8] == 0)
assert np.all(res_nu[:8] == 0)
print("match_all_sync: one lane differs -> pred=0 for every lane in the group")

# 8. Edge case: reading from a lane not in the mask is undefined -> surfaces as NaN
sparse_mask = mask_from_lanes([0, 1, 2, 5, 6])   # lanes 3, 4 are gaps (not active)
sparse_values = np.array([10, 20, 30, 999, 999, 60, 70] + [0] * 25, dtype=np.float64)
down_sparse = shfl_down_sync(sparse_mask, sparse_values, 1)
assert np.isnan(down_sparse[2]), "expected UB(NaN) reading an inactive source lane"
assert down_sparse[0] == 20
print("sparse-mask shuffle: reading an inactive source lane surfaces as NaN (UB)")

# 9. Adversarial: correct partial-warp reduction (13 active lanes, non-power-of-two)
active13 = list(range(13))
vals13 = np.arange(1, 33, dtype=np.float64)
correct_result, used_mask = warp_reduce_sum_partial_correct(vals13, active13)
expected13sum = slow_reduce(vals13, active13)
assert np.isclose(correct_result[active13[0]], expected13sum), (correct_result[active13[0]], expected13sum)
assert used_mask == mask_from_lanes(active13)
print(f"correct partial-warp (13 lanes) reduction: {correct_result[0]} == {expected13sum}")

# 10. Adversarial: same 13 active lanes, but the reduction bug hardcodes FULL_MASK
# and runs the unconditional full-warp loop. This must NOT reproduce the correct
# sum, otherwise the adversarial case proves nothing.
buggy_result = warp_reduce_sum_buggy_full_mask(vals13, active13)
assert not np.isclose(buggy_result, expected13sum), \
    "adversarial case failed to demonstrate corruption: buggy result matched the correct sum"
print(f"buggy full-mask reduction on a 13-lane partial warp: {buggy_result} "
      f"!= correct {expected13sum} (corrupted by stale/garbage inactive-lane data)")

print("\nALL ASSERTS PASSED")

Terminal output from the executed run:

full-warp reduce: lane0=528.0, expected=528.0
butterfly (xor) reduce: all lanes hold 528.0 expected 528.0
shfl_sync broadcast of lane 5: 6.0 6.0
shfl_up_sync delta=1 ok
vote functions: all_sync=0, any_sync=1, ballot matches expected bitmask
all_sync=1 when every active lane's predicate is true
partial-mask (13 lanes) ballot_sync correctly scoped, no bit leakage
match_any_sync groups lanes by value correctly
match_all_sync: uniform group -> pred=1, result=mask
match_all_sync: one lane differs -> pred=0 for every lane in the group
sparse-mask shuffle: reading an inactive source lane surfaces as NaN (UB)
correct partial-warp (13 lanes) reduction: 91.0 == 91.0
buggy full-mask reduction on a 13-lane partial warp: 3075.517959679896 != correct 91.0 (corrupted by stale/garbage inactive-lane data)

ALL ASSERTS PASSED

Failure modes

  • Stale or hardcoded 0xFFFFFFFF mask under divergence. A predicate branches the warp, then a _sync intrinsic inside the branch is called with 0xFFFFFFFF (or a mask computed before the branch) instead of the mask of lanes that actually took that path. Lanes that diverged away never reach the call, so the mask names participants that are not there: undefined behavior. On real hardware this shows up as either a hang (a named lane the hardware waits on never arrives) or a silently wrong value, not a compile error or a clean crash. The "WRONG" comment in the divergence reference template above is this exact bug.
  • A mask whose named lanes do not all reach the same call with the same mask value. This is the general form of the contract the CUDA Programming Guide states for every warp shuffle/vote/match function: even outside of branches, two call sites that both claim to use the "same" mask but compute it differently (e.g. one recomputes __activemask() after a lane has already exited) violate the requirement and are undefined.1
  • Assuming a full warp when only some lanes hold valid data. Distinct from divergence: the whole warp can execute the same instruction uniformly while some lanes carry garbage (a tail iteration, an over-read past an array bound) rather than having taken a different branch. Naming those lanes in the mask does not cause a hang, but it folds garbage into a vote, match, or shuffle result. The adversarial case above reproduces this: a 13-lane partial-warp reduction run with FULL_MASK returns roughly 3075 instead of the correct 91, a silent, non-crashing corruption.
  • Misusing __activemask() as a stable snapshot. __activemask() returns the lanes active at that instruction; it is not guaranteed that the same set of lanes remains convergent at a later instruction unless a _sync intrinsic (including __syncwarp) is used to re-establish it. Treating an old __activemask() result as valid several instructions later reintroduces the stale-mask bug above.
  • Using the deprecated non-_sync intrinsics (__shfl, __shfl_up, __shfl_down, __shfl_xor, __ballot, __any, __all). These take no mask and assume the pre-Volta model where a warp's 32 lanes always move together; they are deprecated and unsafe to rely on for correctness under independent thread scheduling. Replace every one with its _sync counterpart and an explicit, freshly computed mask.

References

Related: GPU Execution Model: SM, Warp, SIMT · CUDA Occupancy Tuning · Warp Specialization & Pipelining · Instruction-Level Parallelism on the GPU · Glossary

  1. NVIDIA CUDA C++ Programming Guide, Warp Shuffle/Vote/Match Functions: the mask argument "indicates the threads participating in the call," "all non-exited threads named in mask must execute the same intrinsic with the same mask, or the result is undefined," and for shuffle specifically, threads may only read data from another thread that is actively participating in the call; if the target thread is inactive, the retrieved value is undefined. ↩↩