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 that a reader must compile; no GPU was available to execute them in this KB's environment. What was independently done here: every device function on this page was assembled into one real, compilable translation unit (warp_primitives_test.cu, reproduced in full in "Executed: compiling and inspecting real device code" below) with __global__ kernel wrappers and a host-side cudaMalloc/launch/cudaDeviceSynchronize harness, compiled with nvcc -arch=sm_70, -arch=sm_80, and -arch=sm_90 (Volta/Ampere/Hopper) with zero errors and zero warnings under -Wall -Wextra, and disassembled with cuobjdump -sass on all three targets to confirm the intrinsics actually lower to the expected hardware instructions (SHFL.DOWN, VOTE.ANY, MATCH.ANY), each instruction attributed to its own mangled kernel symbol using cuobjdump's own function headers, not just that the source parses. A second, independent tool, nvdisasm against the extracted cubin, reproduced byte-identical instruction encodings at the same code offsets. Running the compiled sm_90 binary against this sandbox's driver-less host produced a correctly handled cudaErrorNoDevice from the first cudaMalloc call, not a crash, verifying the host-side error path; compute-sanitizer --tool memcheck was invoked against that same binary and failed for an environment reason (Unable to find injection library libsanitizer-collection.so), documented honestly below rather than omitted. None of this exercises real warp scheduling, divergence, or memory ordering on hardware; see "Executed: compiling and inspecting real device code" for the full transcript and its limits. The numpy block in How to develop with it 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 executing one SIMT instruction stream"]
    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. Every non-exited lane named in mask must execute the same intrinsic with the same mask. Computing a ballot from a branch predicate before divergence is the standard way to form that mask; the invalid case is a mask that names a lane that does not execute the call.

How to use it

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).

// Assembled unmodified into the compiled translation unit in "Executed:
// compiling and inspecting real device code" below. GPU execution needs a device.
__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 safe pattern computes a participation mask from the branch predicate before the warp diverges. Every lane named by that ballot must then execute the intrinsic:

// Assembled unmodified into the compiled translation unit in "Executed:
// compiling and inspecting real device code" below. GPU execution needs a device.
__device__ void process_if_valid(int* data, int idx, int n) {
    bool valid = idx < n;
    unsigned active = __activemask();
    unsigned valid_mask = __ballot_sync(active, valid);

    if (valid) {
        int v = data[idx];
        int peer = __shfl_down_sync(valid_mask, v, 1, 32);
        int lane = threadIdx.x & 31;

        // A source lane outside valid_mask yields an undefined value. Test the
        // source bit before consuming peer; do not infer validity from lane count.
        if (lane < 31 && ((valid_mask >> (lane + 1)) & 1u)) {
            v += peer;
        }
        data[idx] = v;
    }
    // WRONG: using 0xFFFFFFFF inside the branch names lanes that never call.
    // The result is undefined; CUDA does not promise a hang or a specific value.
}

Reference: vote, match, and syncwarp

// Assembled unmodified into the compiled translation unit in "Executed:
// compiling and inspecting real device code" below. GPU execution needs a device.
__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;
}

Executed: compiling and inspecting real device code

An earlier revision of this page described a full translation unit and pasted a SASS transcript for it, but the repository did not actually contain that file, and the pasted transcript covered only four of the five reduction offsets while claiming to cover all five. This section replaces that with a real file, compiled and disassembled in this session, with commands and output pasted verbatim.

The four device functions above (warp_reduce_sum, process_if_valid, warp_aggregated_histogram, warp_any_out_of_bounds) were assembled, unmodified, into one real .cu translation unit, warp_primitives_test.cu, reproduced here in full:

// warp_primitives_test.cu
//
// Full, compilable translation unit assembling the device functions described
// in docs/warp-level-cuda-primitives.md: warp_reduce_sum (shuffle-down
// butterfly reduction), process_if_valid (divergence-safe shuffle under a
// computed mask), warp_aggregated_histogram (match + vote + syncwarp), and
// warp_any_out_of_bounds (any_sync). Each device function gets a thin
// __global__ kernel wrapper. host main() allocates device memory, initializes
// every device input buffer from a host array (d_keys is clamped mod NUM_BINS
// so its values never index past the bins array actually allocated), launches
// every kernel, and checks every CUDA API return code. This file was compiled
// with nvcc for sm_70, sm_80, and sm_90 and disassembled with cuobjdump in a
// sandbox with no bound GPU; it was not executed against a device.

#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <cuda_runtime.h>

#define CUDA_CHECK(call)                                                     \
    do {                                                                     \
        cudaError_t _err = (call);                                           \
        if (_err != cudaSuccess) {                                           \
            fprintf(stderr, "CUDA error at %s:%d: %s\n", __FILE__, __LINE__, \
                    cudaGetErrorString(_err));                               \
            return 1;                                                        \
        }                                                                    \
    } while (0)

// ---------------------------------------------------------------------------
// Device functions (unmodified from the KB page)
// ---------------------------------------------------------------------------

__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
}

__device__ void process_if_valid(int* data, int idx, int n) {
    bool valid = idx < n;
    unsigned active = __activemask();
    unsigned valid_mask = __ballot_sync(active, valid);

    if (valid) {
        int v = data[idx];
        int peer = __shfl_down_sync(valid_mask, v, 1, 32);
        int lane = threadIdx.x & 31;

        // A source lane outside valid_mask yields an undefined value. Test the
        // source bit before consuming peer; do not infer validity from lane count.
        if (lane < 31 && ((valid_mask >> (lane + 1)) & 1u)) {
            v += peer;
        }
        data[idx] = v;
    }
    // WRONG: using 0xFFFFFFFF inside the branch names lanes that never call.
    // The result is undefined; CUDA does not promise a hang or a specific value.
}

__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;
}

// ---------------------------------------------------------------------------
// __global__ kernel wrappers, one per device function
// ---------------------------------------------------------------------------

__global__ void kernel_warp_reduce_sum(const float* in, float* out, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    float v = (idx < n) ? in[idx] : 0.0f;
    float sum = warp_reduce_sum(v);
    if ((threadIdx.x & 31) == 0) out[idx / 32] = sum;
}

__global__ void kernel_process_if_valid(int* data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    process_if_valid(data, idx, n);
}

__global__ void kernel_warp_aggregated_histogram(int* bins, const int* keys, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        warp_aggregated_histogram(bins, keys[idx]);
    }
}

__global__ void kernel_warp_any_out_of_bounds(const int* idxs, int n, int bound, bool* out) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {
        bool r = warp_any_out_of_bounds(idxs[tid], bound);
        if ((threadIdx.x & 31) == 0) out[tid / 32] = r;
    }
}

// Deliberately undefined-behavior-by-construction kernel: 0xFFFFFFFF names
// all 32 lanes, but lanes with idx >= n never reach the __shfl_down_sync
// call, so those named lanes do not execute the same intrinsic with the same
// mask. Never launched by default; gated behind --invalid-mask so only a
// deliberate Compute Sanitizer pass (on a real GPU) launches it.
__global__ void k_invalid_mask(int* data, int n) {
    int idx = threadIdx.x;
    if (idx < n) {
        int peer = __shfl_down_sync(0xFFFFFFFF, data[idx], 1);
        data[idx] += peer;
    }
}

// ---------------------------------------------------------------------------
// Host harness
// ---------------------------------------------------------------------------

int main(int argc, char** argv) {
    bool run_invalid_mask = false;
    for (int i = 1; i < argc; ++i) {
        if (0 == strcmp(argv[i], "--invalid-mask")) run_invalid_mask = true;
    }

    const int N = 1024;
    const int NUM_WARPS = N / 32;
    const int NUM_BINS = 16;

    float *d_in = nullptr, *d_out = nullptr;
    cudaError_t err = cudaMalloc(&d_in, N * sizeof(float));
    if (err != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed (expected on a host with no bound "
                         "GPU): %s\n", cudaGetErrorString(err));
        return 2;   // fail closed: compilation succeeded, device validation did not run
    }

    CUDA_CHECK(cudaMalloc(&d_out, NUM_WARPS * sizeof(float)));

    int *d_data = nullptr, *d_bins = nullptr, *d_keys = nullptr;
    bool *d_oob_out = nullptr;
    CUDA_CHECK(cudaMalloc(&d_data, N * sizeof(int)));
    CUDA_CHECK(cudaMalloc(&d_bins, NUM_BINS * sizeof(int)));
    CUDA_CHECK(cudaMalloc(&d_keys, N * sizeof(int)));
    CUDA_CHECK(cudaMalloc(&d_oob_out, NUM_WARPS * sizeof(bool)));

    // cudaMalloc does not zero or otherwise initialize device memory: every
    // buffer a kernel reads must be populated from a known host array first.
    // d_keys in particular drives atomicAdd(&bins[key], ...) inside
    // warp_aggregated_histogram; leaving it as uninitialized garbage would
    // let an arbitrary int32 value index far outside the NUM_BINS-sized bins
    // array and issue an out-of-bounds atomic on real hardware. Keys are
    // therefore clamped (mod NUM_BINS) to the exact bin count allocated
    // above, and d_in/d_data get deterministic values instead of whatever
    // bytes cudaMalloc happened to return.
    float* h_in = static_cast<float*>(malloc(N * sizeof(float)));
    int* h_data = static_cast<int*>(malloc(N * sizeof(int)));
    int* h_keys = static_cast<int*>(malloc(N * sizeof(int)));
    for (int i = 0; i < N; ++i) {
        h_in[i] = static_cast<float>(i + 1);
        h_data[i] = i;
        h_keys[i] = i % NUM_BINS;
    }
    CUDA_CHECK(cudaMemcpy(d_in, h_in, N * sizeof(float), cudaMemcpyHostToDevice));
    CUDA_CHECK(cudaMemcpy(d_data, h_data, N * sizeof(int), cudaMemcpyHostToDevice));
    CUDA_CHECK(cudaMemcpy(d_keys, h_keys, N * sizeof(int), cudaMemcpyHostToDevice));
    free(h_in);
    free(h_data);
    free(h_keys);

    dim3 block(128);
    dim3 grid((N + block.x - 1) / block.x);

    kernel_warp_reduce_sum<<<grid, block>>>(d_in, d_out, N);
    CUDA_CHECK(cudaGetLastError());

    kernel_process_if_valid<<<grid, block>>>(d_data, N - 5);
    CUDA_CHECK(cudaGetLastError());

    CUDA_CHECK(cudaMemset(d_bins, 0, NUM_BINS * sizeof(int)));
    kernel_warp_aggregated_histogram<<<grid, block>>>(d_bins, d_keys, N);
    CUDA_CHECK(cudaGetLastError());

    kernel_warp_any_out_of_bounds<<<grid, block>>>(d_keys, N, N - 5, d_oob_out);
    CUDA_CHECK(cudaGetLastError());

    if (run_invalid_mask) {
        k_invalid_mask<<<1, 32>>>(d_data, 17);
        CUDA_CHECK(cudaGetLastError());
    }

    CUDA_CHECK(cudaDeviceSynchronize());

    cudaFree(d_in);
    cudaFree(d_out);
    cudaFree(d_data);
    cudaFree(d_bins);
    cudaFree(d_keys);
    cudaFree(d_oob_out);

    printf("All kernels launched and synchronized successfully.\n");
    return 0;
}

Correction applied in this revision. An earlier revision of this harness called cudaMalloc on d_in, d_data, and d_keys and then launched kernels directly, without ever writing to those buffers first. d_keys fed atomicAdd(&bins[key], count) with whatever uninitialized bytes cudaMalloc returned, so on a real GPU an arbitrary key value could drive an out-of-bounds atomic write past the 16-entry bins array; d_in and d_data were likewise read by their kernels while holding garbage rather than known values. The main() above now populates all three from host arrays via cudaMemcpy before any kernel launch, with d_keys explicitly clamped mod NUM_BINS to match the bin count actually allocated. This is a source-level fix, checked only by recompiling (below) and by re-reading the corrected code; it has not been run on a GPU, so it does not by itself demonstrate the absence of out-of-bounds atomics at runtime, only that the indices supplied now are constructed to stay in range. Device execution, a Compute Sanitizer pass, and any numerical-equivalence, divergence, or memory-ordering check remain unexecuted in this environment, exactly as the rest of this section already documents.

Compiler and toolchain identity, then all three architecture targets, compiled with warnings enabled on the host compiler:

$ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2023 NVIDIA Corporation
Built on Fri_Jan__6_16:45:21_PST_2023
Cuda compilation tools, release 12.0, V12.0.140
Build cuda_12.0.r12.0/compiler.32267302_0

$ nvcc -arch=sm_70 -Xcompiler -Wall -Xcompiler -Wextra -o warp_test_sm_70 warp_primitives_test.cu
$ echo "sm_70 exit: $?"
sm_70 exit: 0

$ nvcc -arch=sm_80 -Xcompiler -Wall -Xcompiler -Wextra -o warp_test_sm_80 warp_primitives_test.cu
$ echo "sm_80 exit: $?"
sm_80 exit: 0

$ nvcc -arch=sm_90 -Xcompiler -Wall -Xcompiler -Wextra -o warp_test_sm_90 warp_primitives_test.cu
$ echo "sm_90 exit: $?"
sm_90 exit: 0

All three architecture targets compiled and linked against libcudart with zero errors and zero warnings, -Wall -Wextra included. A fourth build, combining all three targets into one fatbin with an explicit host compiler, also succeeded:

$ nvcc -ccbin g++-12 \
    -gencode arch=compute_70,code=sm_70 \
    -gencode arch=compute_80,code=sm_80 \
    -gencode arch=compute_90,code=sm_90 \
    -Xcompiler -Wall -Xcompiler -Wextra \
    -o warp_test_fatbin warp_primitives_test.cu
$ echo $?
0
$ cuobjdump --list-elf warp_test_fatbin
ELF file    1: tmpxft_00012aae_00000000-0.sm_70.cubin
ELF file    2: tmpxft_00012aae_00000000-1.sm_80.cubin
ELF file    3: tmpxft_00012aae_00000000-2.sm_90.cubin
ELF file    4: tmpxft_00012aae_00000000-3.sm_70.cubin
ELF file    5: tmpxft_00012aae_00000000-4.sm_80.cubin
ELF file    6: tmpxft_00012aae_00000000-5.sm_90.cubin

Running the sm_90 binary in this sandbox (no bound NVIDIA driver):

$ ./warp_test_sm_90
cudaMalloc failed (expected on a host with no bound GPU): no CUDA-capable device is detected
$ echo $?
2

The real cudaMalloc call failed exactly as cudaGetErrorString documents (cudaErrorNoDevice), and the host code's own error path caught it without segfaulting or hanging. It returns nonzero because a no-device run did not validate any kernel; an earlier revision returned 0 here, which allowed CI to misclassify a compile-only result as a successful device test.

SASS disassembly, attributed per kernel. cuobjdump -sass on the compiled sm_90 binary lists five function symbols, one per kernel in the translation unit above:

$ cuobjdump -sass warp_test_sm_90 > full_sass_sm90.txt
$ grep -n "Function :" full_sass_sm90.txt
19:     Function : _Z14k_invalid_maskPii
72:     Function : _Z29kernel_warp_any_out_of_boundsPKiiiPb
157:        Function : _Z32kernel_warp_aggregated_histogramPiPKii
242:        Function : _Z23kernel_process_if_validPii
327:        Function : _Z22kernel_warp_reduce_sumPKfPfi

Demangled with c++filt: k_invalid_mask(int*, int), kernel_warp_any_out_of_bounds(int const*, int, int, bool*), kernel_warp_aggregated_histogram(int*, int const*, int), kernel_process_if_valid(int*, int), kernel_warp_reduce_sum(float const*, float*, int).

kernel_warp_reduce_sum is the last function in the file, so its instructions run from line 327 to the next .......... padding marker. Scoping the listing to that one function's address range, before grepping for SHFL.DOWN, is what keeps the five reduction instructions from being mixed with any other kernel's shuffle:

$ awk '/Function : _Z22kernel_warp_reduce_sum/{f=1} f{print} f && /\.\.\.\.\.\.\.\.\.\./{exit}' full_sass_sm90.txt | grep "SHFL.DOWN"
        /*00e0*/                   SHFL.DOWN PT, R3, R2, 0x10, 0x1f ;            /* 0x0a001f0002037f89 */
        /*0110*/                   SHFL.DOWN PT, R0, R3, 0x8, 0x1f ;             /* 0x09001f0003007f89 */
        /*0130*/                   SHFL.DOWN PT, R7, R0, 0x4, 0x1f ;             /* 0x08801f0000077f89 */
        /*0150*/                   SHFL.DOWN PT, R4, R7, 0x2, 0x1f ;             /* 0x08401f0007047f89 */
        /*0170*/                   SHFL.DOWN PT, R9, R4, 0x1, 0x1f ;             /* 0x08201f0004097f89 */

All five offsets, 0x10, 0x8, 0x4, 0x2, 0x1 (16, 8, 4, 2, 1 in decimal), are present, in program order, at addresses /*00e0*/ through /*0170*/, entirely inside kernel_warp_reduce_sum's own function body (line 327 onward). This is the complete halving sequence warp_reduce_sum's C++ loop specifies, for (offset = 16; offset > 0; offset >>= 1), one shuffle per iteration, with 0x2 present.

The other four kernels each contribute their own, separately attributed instructions from the same full_sass_sm90.txt, none merged with the reduction kernel's sequence:

  • k_invalid_mask (function header at line 19): one SHFL.DOWN PT, R5, R0, 0x1, 0x1f at /*0090*/, its single unconditional __shfl_down_sync(0xFFFFFFFF, data[idx], 1) call.
  • kernel_warp_any_out_of_bounds (line 72): two VOTE.ANY instructions, VOTE.ANY R0, PT, PT at /*00e0*/ and VOTE.ANY P0, P0 at /*0110*/, from __activemask() and __any_sync.
  • kernel_warp_aggregated_histogram (line 157): VOTE.ANY R0, PT, PT at /*00c0*/ from __activemask(), MATCH.ANY R4, R5 at /*0100*/ from __match_any_sync, and the atomic add lowers to REDG.E.ADD.STRONG.GPU at /*0180*/ rather than a discrete ATOM instruction, because the return value of atomicAdd is unused and the compiler folds it into a reduction.
  • kernel_process_if_valid (line 242): two VOTE.ANY instructions from __activemask()/__ballot_sync, plus one, unrelated SHFL.DOWN PT, R4, R5, 0x1, 0x1f at /*0150*/, from its own single __shfl_down_sync(valid_mask, v, 1, 32) call. Conflating this line with kernel_warp_reduce_sum's five-instruction sequence, by grepping the whole binary instead of scoping to one function's address range, was the exact defect in an earlier revision of this page; the function-scoped listing above is how that is avoided this time.

As a second, independent disassembler, the sm_90 cubin was extracted from the binary and run through nvdisasm:

$ cuobjdump --extract-elf all warp_test_sm_90
Extracting ELF file    1: tmpxft_00009c5f_00000000-0.sm_90.cubin
Extracting ELF file    2: tmpxft_00009c5f_00000000-1.sm_90.cubin
$ nvdisasm -c tmpxft_00009c5f_00000000-1.sm_90.cubin | grep "SHFL.DOWN"
        /*0090*/                   SHFL.DOWN PT, R5, R0, 0x1, 0x1f ;
        /*0150*/                   SHFL.DOWN PT, R4, R5, 0x1, 0x1f ;
        /*00e0*/                   SHFL.DOWN PT, R3, R2, 0x10, 0x1f ;
        /*0110*/                   SHFL.DOWN PT, R0, R3, 0x8, 0x1f ;
        /*0130*/                   SHFL.DOWN PT, R7, R0, 0x4, 0x1f ;
        /*0150*/                   SHFL.DOWN PT, R4, R7, 0x2, 0x1f ;
        /*0170*/                   SHFL.DOWN PT, R9, R4, 0x1, 0x1f ;

nvdisasm reproduces the same instruction text and the same code offsets as cuobjdump -sass for every SHFL.DOWN in the binary (nvdisasm -c does not print function-symbol headers, which is why the cuobjdump listing above, not this one, is the source for kernel attribution; the two tools are cross-checked against each other on instruction content, not on symbol names).

sm_70 and sm_80 were checked the same way, scoping each dump to kernel_warp_reduce_sum's own function range. Both produce the identical five-offset sequence, at different addresses because the older architecture targets emit slightly different setup code ahead of the first shuffle:

$ awk '/Function : _Z22kernel_warp_reduce_sum/{f=1} f{print} f && /\.\.\.\.\.\.\.\.\.\./{exit}' full_sass_sm70.txt | grep "SHFL.DOWN"
        /*00e0*/                   SHFL.DOWN PT, R3, R2, 0x10, 0x1f ;
        /*0110*/                   SHFL.DOWN PT, R4, R3, 0x8, 0x1f ;
        /*0130*/                   SHFL.DOWN PT, R5, R4, 0x4, 0x1f ;
        /*0150*/                   SHFL.DOWN PT, R6, R5, 0x2, 0x1f ;
        /*0170*/                   SHFL.DOWN PT, R7, R6, 0x1, 0x1f ;

$ awk '/Function : _Z22kernel_warp_reduce_sum/{f=1} f{print} f && /\.\.\.\.\.\.\.\.\.\./{exit}' full_sass_sm80.txt | grep "SHFL.DOWN"
        /*00d0*/                   SHFL.DOWN PT, R3, R2, 0x10, 0x1f ;
        /*0100*/                   SHFL.DOWN PT, R4, R3, 0x8, 0x1f ;
        /*0120*/                   SHFL.DOWN PT, R5, R4, 0x4, 0x1f ;
        /*0140*/                   SHFL.DOWN PT, R6, R5, 0x2, 0x1f ;
        /*0160*/                   SHFL.DOWN PT, R7, R6, 0x1, 0x1f ;

What was not verified, honestly. compute-sanitizer --tool memcheck was invoked against the compiled sm_90 binary in this session and failed for an environment reason, not a code defect:

$ compute-sanitizer --tool memcheck ./warp_test_sm_90
========= COMPUTE-SANITIZER
========= Unable to find injection library libsanitizer-collection.so
$ echo $?
13

This sandbox's CUDA toolkit install does not carry a complete Compute Sanitizer runtime, and there is no GPU to run a sanitizer pass against regardless. The binary itself was never launched against a device: k_invalid_mask compiled cleanly as the fifth kernel above (see its own SHFL.DOWN and function header at line 19) and is gated behind the --invalid-mask argument so it never runs in the default path, but no run of either the default path or --invalid-mask reached a GPU in this session. This translation unit also does not implement a CPU-reference equivalence check inside main(); adding one would still not exercise real hardware without a device, so it is left as a documented gap rather than claimed as done. Nothing on this page has verified real warp scheduling, actual divergence behavior, memory ordering under __syncwarp, or numerical correctness on hardware. That requires the binary above run with compute-sanitizer --tool racecheck, memcheck, and synccheck on a real Volta-or-later GPU, both on the default path and with --invalid-mask (the latter must produce a true-positive flag from synccheck, not silence), exactly the CPU-equivalence-plus-Sanitizer validation the numpy model below cannot substitute for.

How to develop with it

Runnable software model

The block below models the data movement of all four intrinsic families and rejects a call when the named and actual participants differ. An inactive shuffle source is represented as NaN because CUDA leaves that value undefined. The model cannot reproduce GPU scheduling or __syncwarp memory ordering, so it tests participation only. It was executed with python3; every assertion and the pasted output were checked.

"""
Software model of warp-level CUDA shuffle, vote, match, and participation.
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 the CUDA intrinsics. The model checks that the
caller set equals the named set. It represents an inactive shuffle source as
NaN. It does not model GPU scheduling, instruction issue, or memory ordering.
"""
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 named_lanes(mask):
    return {lane for lane in range(WARP_SIZE) if lane_active(mask, lane)}


def validate_participants(mask, callers):
    named = named_lanes(mask)
    actual = set(callers)
    if named != actual:
        raise RuntimeError(
            f"participant mismatch: named={sorted(named)}, callers={sorted(actual)}"
        )


def syncwarp(mask, callers):
    """Check participation only; this model cannot reproduce memory ordering."""
    validate_participants(mask, callers)


def shfl_down_sync(mask, values, delta, width=WARP_SIZE, callers=None):
    """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."""
    if callers is not None:
        validate_participants(mask, callers)
    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):
    """Reduce a contiguous, 0-based partial warp with a valid ballot mask."""
    n = len(active_lanes)
    mask = mask_from_lanes(active_lanes)
    v = values.astype(np.float64).copy()
    offset = 1
    while offset < n:
        shuffled = shfl_down_sync(mask, v, offset, callers=active_lanes)
        for i, lane in enumerate(active_lanes):
            if i + offset < n:
                v[lane] += shuffled[lane]
        offset *= 2
    return v, mask


# ---------------------------------------------------------------------------
# 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")

syncwarp(partial_mask, partial_lanes)
print("syncwarp participation contract: 13 named lanes arrived")


# 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: FULL_MASK names 19 lanes that do not execute the call.
try:
    shfl_down_sync(FULL_MASK, vals13, 16, callers=active13)
except RuntimeError as exc:
    assert "participant mismatch" in str(exc)
else:
    raise AssertionError("invalid full-mask call was not rejected")
print("invalid full-mask call: missing participants rejected by model")

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
syncwarp participation contract: 13 named lanes arrived
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
invalid full-mask call: missing participants rejected by model

ALL ASSERTS PASSED

How to maintain it

Keep the participation predicate beside the ballot that creates its mask. In code review, trace every named lane to the same intrinsic call and reject a later __activemask() that silently narrows the intended group. Compile CUDA reference tests for each supported architecture, the same nvcc -arch=sm_XX plus cuobjdump -sass check exercised above, so a compiler or CUDA-version upgrade that changes instruction selection is visible in review, and run Compute Sanitizer (memcheck and racecheck) on boundary-sized inputs on real hardware, something this page's own environment could not do (see the executed section above for the exact "missing injection library" failure it hit).

How to run it in production

Compare each optimized kernel against a CPU reference for sizes 0, 1, 31, 32, and 33 plus sparse predicates. Guard both call participation and shuffle-source validity at tail boundaries. Profile eligible warps and branch efficiency after compiler or CUDA upgrades; the Python model does not validate generated instructions, scheduling, or fences, and neither does a clean nvcc compile by itself, disassemble the compiled binary (cuobjdump -sass) after any compiler or architecture-target change and diff the shuffle/vote/match instruction sequence against a known-good baseline, the way this page's own executed section does, before trusting that source-level correctness survived the recompile.

Failure modes

  • A mask names lanes that do not execute the call. A ballot computed from the branch predicate before divergence is valid. A hardcoded full mask, or a stale mask that includes lanes that exited or took another branch, violates the participation contract. The result is undefined; CUDA does not promise a hang, a trap, or a stable wrong value.
  • 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
  • Confusing participation with data validity. If all named lanes execute but some loaded out-of-bounds or stale data, the intrinsic can satisfy its mask contract while the algorithm remains wrong. Bounds checks and source-lane guards are separate from participation checks; the runnable model does not conflate the two.
  • 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. ↩↩