Warp-Level Programming

Exchange registers between threads of a warp without shared memory:

13.1 Warp shuffle

Exchange registers between threads of a warp without shared memory:

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

__inline__ __device__ float warpReduceSum(float val) {
    for (int offset = warpSize / 2; offset > 0; offset >>= 1)
        val += __shfl_down_sync(0xffffffff, val, offset);
    return val;                                    // lane 0 holds the warp's sum
}

__inline__ __device__ float warpAllReduceSum(float val) {
    for (int mask = 16; mask > 0; mask >>= 1)
        val += __shfl_xor_sync(0xffffffff, val, mask);   // butterfly: every lane gets the sum
    return val;
}

13.2 Warp vote & match functions

Example — warp-aggregated atomic increment:

int  __all_sync(mask, pred);       // true if pred is true for all active threads
int  __any_sync(mask, pred);
unsigned __ballot_sync(mask, pred);// bitmask of threads where pred is true
unsigned __activemask();
unsigned __match_any_sync(mask, value);   // threads with equal values (CC 7.0+)
int lane = threadIdx.x % 32;  // or threadIdx.x & 31
int warpId = threadIdx.x / 32;
__popc(ballot);                    // count set bits
__ffs(x);                          // find first set bit

__device__ int atomicAggInc(int* ctr) {
    unsigned active = __activemask();
    int leader = __ffs(active) - 1;
    int lane = threadIdx.x & 31;
    int res;
    if (lane == leader) res = atomicAdd(ctr, __popc(active));
    res = __shfl_sync(active, res, leader);
    return res + __popc(active & ((1u << lane) - 1));
}

13.3 Warp divergence

Minimize divergence in hot loops; arrange data so warps take uniform paths.

if (threadIdx.x % 2 == 0) a(); else b();       // divergent: both paths executed serially
if ((threadIdx.x / 32) % 2 == 0) a(); else b(); // warp-uniform: no divergence

13.4 Warp-level matrix ops & reductions (CC 8.0+)

__reduce_add_sync(mask, value), __reduce_min_sync, __reduce_max_sync for integers; WMMA/MMA for Tensor Cores (§27).