Warp-level primitives

Lanes 0..31 of a warp are consecutive linear thread ids (tid % 32 == laneId). Shuffle/vote/reduce are collective: every participating lane deposits a value, rendezvous at a warp-wide barrier (the scheduler releases the warp only once all its live lanes have arrived), then reads peers from the snapshot. Values travel as an 8-byte blob so any <=8-byte type (int, unsigned, float, long long, double) shuffles correctly.

Defined in code/simulator.h.

This page documents 15 symbols: 1 constants, 14 functions.

Constants

Signature	Description	Availability	Location	Example
inline int warpSize = 32	Lanes 0..31 of a warp are consecutive linear thread ids (tid % 32 == laneId).	sim only	simulator.h:1007

Functions

Signature	Description	Availability	Location	Example
inline unsigned __activemask()	__activemask does NOT synchronise; approximate the converged set as every real lane of the warp (good enough for the common ballot/shuffle feeding uses).	sim only	simulator.h:1103
inline int __all_sync(unsigned mask, int pred)	Non-zero iff the predicate holds for EVERY participating lane of the warp.	sim only	simulator.h:1093	View
inline int __any_sync(unsigned mask, int pred)	Non-zero iff the predicate holds for ANY participating lane of the warp.	sim only	simulator.h:1084	View
inline unsigned __ballot_sync(unsigned mask, int pred)	Vote: ballot returns the bitmask of participating lanes whose predicate holds.	sim only	simulator.h:1074	View
inline unsigned __reduce_add_sync(unsigned mask, unsigned val)	Warp reductions. NOTE: the real intrinsics need sm_80 (Ampere) — they do NOT exist on this repo’s sm_75 target, so a portable kernel can’t use them on a T4. Provided here for sim completeness / …	sim only	simulator.h:1114
inline unsigned __reduce_max_sync(unsigned mask, unsigned val)		sim only	simulator.h:1134
inline unsigned __reduce_min_sync(unsigned mask, unsigned val)		sim only	simulator.h:1123
inline T __shfl_down_sync(unsigned mask, T var, unsigned delta, int width = 32)	Shift down: read lane (self+delta); lanes past the subsection end keep their own.	sim only	simulator.h:1039	View
inline T __shfl_sync(unsigned mask, T var, int srcLane, int width = 32)	Indexed broadcast: read lane srcLane within the caller’s width-subsection.	sim only	simulator.h:1018	View
inline T __shfl_up_sync(unsigned mask, T var, unsigned delta, int width = 32)	Shift up: read lane (self-delta); lanes that would read before the subsection base keep their own value.	sim only	simulator.h:1029	View
inline T __shfl_xor_sync(unsigned mask, T var, int laneMask, int width = 32)	Butterfly: read lane (self ^ laneMask) within the subsection.	sim only	simulator.h:1050	View
inline void __syncwarp(unsigned mask = 0xFFFFFFFFu)	__syncwarp: warp-wide barrier with no data exchange.	sim only	simulator.h:1062	View
inline T sim_dec(uint64_t u)		sim only	simulator.h:1013
inline uint64_t sim_enc(T v)	Pack/unpack a shuffle payload into the 8-byte exchange slot (any <=8-byte type).	sim only	simulator.h:1010

Examples

Code from the repository’s example corpus (code\*.cu); the highlighted line is the call site. These blocks are what the View links above open.

__all_sync (test_warp.cu:74)

    for (int off = 16; off > 0; off >>= 1) v += __shfl_xor_sync(FULL, v, off);
    out[gid] = v;  // every lane ends with the full warp sum
}

__global__ void kBallot(const float* in, unsigned* outWarp, int* outAny, int* outAll) {
    const unsigned gid  = blockIdx.x * blockDim.x + threadIdx.x;
    const unsigned lane = threadIdx.x % 32u;
    const int pred = in[gid] > 6.0f;
    const unsigned b = __ballot_sync(FULL, pred);
    const int any = __any_sync(FULL, pred);
    const int all = __all_sync(FULL, pred);
    if (lane == 0) { outWarp[gid / 32u] = b; outAny[gid / 32u] = any; outAll[gid / 32u] = all; }
}

// Partial-warp ballot: a single block of `n` threads (n not a multiple of 32).
__global__ void kBallotPartial(unsigned* outWarp) {
    const unsigned tid  = threadIdx.x;
    const unsigned lane = tid % 32u;
    const unsigned b = __ballot_sync(FULL, 1);  // predicate always true
    if (lane == 0) outWarp[tid / 32u] = b;      // -> mask of ACTIVE lanes only
}

__any_sync (test_warp.cu:73)

    float v = in[gid];
    for (int off = 16; off > 0; off >>= 1) v += __shfl_xor_sync(FULL, v, off);
    out[gid] = v;  // every lane ends with the full warp sum
}

__global__ void kBallot(const float* in, unsigned* outWarp, int* outAny, int* outAll) {
    const unsigned gid  = blockIdx.x * blockDim.x + threadIdx.x;
    const unsigned lane = threadIdx.x % 32u;
    const int pred = in[gid] > 6.0f;
    const unsigned b = __ballot_sync(FULL, pred);
    const int any = __any_sync(FULL, pred);
    const int all = __all_sync(FULL, pred);
    if (lane == 0) { outWarp[gid / 32u] = b; outAny[gid / 32u] = any; outAll[gid / 32u] = all; }
}

// Partial-warp ballot: a single block of `n` threads (n not a multiple of 32).
__global__ void kBallotPartial(unsigned* outWarp) {
    const unsigned tid  = threadIdx.x;
    const unsigned lane = tid % 32u;
    const unsigned b = __ballot_sync(FULL, 1);  // predicate always true
    if (lane == 0) outWarp[tid / 32u] = b;      // -> mask of ACTIVE lanes only
}

__ballot_sync (test_warp.cu:72)

    const unsigned gid = blockIdx.x * blockDim.x + threadIdx.x;
    float v = in[gid];
    for (int off = 16; off > 0; off >>= 1) v += __shfl_xor_sync(FULL, v, off);
    out[gid] = v;  // every lane ends with the full warp sum
}

__global__ void kBallot(const float* in, unsigned* outWarp, int* outAny, int* outAll) {
    const unsigned gid  = blockIdx.x * blockDim.x + threadIdx.x;
    const unsigned lane = threadIdx.x % 32u;
    const int pred = in[gid] > 6.0f;
    const unsigned b = __ballot_sync(FULL, pred);
    const int any = __any_sync(FULL, pred);
    const int all = __all_sync(FULL, pred);
    if (lane == 0) { outWarp[gid / 32u] = b; outAny[gid / 32u] = any; outAll[gid / 32u] = all; }
}

// Partial-warp ballot: a single block of `n` threads (n not a multiple of 32).
__global__ void kBallotPartial(unsigned* outWarp) {
    const unsigned tid  = threadIdx.x;
    const unsigned lane = tid % 32u;
    const unsigned b = __ballot_sync(FULL, 1);  // predicate always true
    if (lane == 0) outWarp[tid / 32u] = b;      // -> mask of ACTIVE lanes only

__shfl_down_sync (test_warp.cu:41)

constexpr int N     = 256;   // 8 warps
constexpr int BLOCK = 128;   // 4 warps/block, multiple of 32
constexpr int WARPS = N / 32;

// ---- Kernels ---------------------------------------------------------------

__global__ void kReduceSum(const float* in, float* outWarp) {
    const unsigned gid  = blockIdx.x * blockDim.x + threadIdx.x;
    const unsigned lane = threadIdx.x % 32u;
    float v = in[gid];
    for (int off = 16; off > 0; off >>= 1) v += __shfl_down_sync(FULL, v, off);
    if (lane == 0) outWarp[gid / 32u] = v;  // lane 0 holds the warp sum
}

__global__ void kBroadcast(const float* in, float* out) {
    const unsigned gid = blockIdx.x * blockDim.x + threadIdx.x;
    out[gid] = __shfl_sync(FULL, in[gid], 0);  // every lane <- lane 0's value
}

__global__ void kInclusiveScan(const float* in, float* out) {
    const unsigned gid  = blockIdx.x * blockDim.x + threadIdx.x;
    const unsigned lane = threadIdx.x % 32u;

__shfl_sync (test_warp.cu:47)

__global__ void kReduceSum(const float* in, float* outWarp) {
    const unsigned gid  = blockIdx.x * blockDim.x + threadIdx.x;
    const unsigned lane = threadIdx.x % 32u;
    float v = in[gid];
    for (int off = 16; off > 0; off >>= 1) v += __shfl_down_sync(FULL, v, off);
    if (lane == 0) outWarp[gid / 32u] = v;  // lane 0 holds the warp sum
}

__global__ void kBroadcast(const float* in, float* out) {
    const unsigned gid = blockIdx.x * blockDim.x + threadIdx.x;
    out[gid] = __shfl_sync(FULL, in[gid], 0);  // every lane <- lane 0's value
}

__global__ void kInclusiveScan(const float* in, float* out) {
    const unsigned gid  = blockIdx.x * blockDim.x + threadIdx.x;
    const unsigned lane = threadIdx.x % 32u;
    float v = in[gid];
    for (int off = 1; off < 32; off <<= 1) {
        float n = __shfl_up_sync(FULL, v, off);
        if (lane >= static_cast<unsigned>(off)) v += n;
    }
    out[gid] = v;

__shfl_up_sync (test_warp.cu:55)

__global__ void kBroadcast(const float* in, float* out) {
    const unsigned gid = blockIdx.x * blockDim.x + threadIdx.x;
    out[gid] = __shfl_sync(FULL, in[gid], 0);  // every lane <- lane 0's value
}

__global__ void kInclusiveScan(const float* in, float* out) {
    const unsigned gid  = blockIdx.x * blockDim.x + threadIdx.x;
    const unsigned lane = threadIdx.x % 32u;
    float v = in[gid];
    for (int off = 1; off < 32; off <<= 1) {
        float n = __shfl_up_sync(FULL, v, off);
        if (lane >= static_cast<unsigned>(off)) v += n;
    }
    out[gid] = v;
}

__global__ void kXorSum(const float* in, float* out) {
    const unsigned gid = blockIdx.x * blockDim.x + threadIdx.x;
    float v = in[gid];
    for (int off = 16; off > 0; off >>= 1) v += __shfl_xor_sync(FULL, v, off);
    out[gid] = v;  // every lane ends with the full warp sum
}

__shfl_xor_sync (test_warp.cu:64)

    for (int off = 1; off < 32; off <<= 1) {
        float n = __shfl_up_sync(FULL, v, off);
        if (lane >= static_cast<unsigned>(off)) v += n;
    }
    out[gid] = v;
}

__global__ void kXorSum(const float* in, float* out) {
    const unsigned gid = blockIdx.x * blockDim.x + threadIdx.x;
    float v = in[gid];
    for (int off = 16; off > 0; off >>= 1) v += __shfl_xor_sync(FULL, v, off);
    out[gid] = v;  // every lane ends with the full warp sum
}

__global__ void kBallot(const float* in, unsigned* outWarp, int* outAny, int* outAll) {
    const unsigned gid  = blockIdx.x * blockDim.x + threadIdx.x;
    const unsigned lane = threadIdx.x % 32u;
    const int pred = in[gid] > 6.0f;
    const unsigned b = __ballot_sync(FULL, pred);
    const int any = __any_sync(FULL, pred);
    const int all = __all_sync(FULL, pred);
    if (lane == 0) { outWarp[gid / 32u] = b; outAny[gid / 32u] = any; outAll[gid / 32u] = all; }

__syncwarp (test_web_bf_warp_cg.cu:71)

        }
        __syncthreads();
    }

    for (int offset = 16; offset > 0; offset >>= 1)
    {
        if (tid < offset)
        {
            s_y[tid] += s_y[tid + offset];
        }
        __syncwarp();
    }

    if (tid == 0)
    {
        atomicAdd(d_y, s_y[0]);
    }
}

// verbatim device kernel — __shfl_down_sync for the last warp
void __global__ reduce_shfl(const real *d_x, real *d_y, const int N)
{