Stage 4 · Going lower11 / 15about 24 mincuda/09_warp_shuffle.cu
Warp primitives: exchanging data without shared memory
Shuffle, ballot, cooperative groups, and how to use atomics correctly
After this lesson you can
- Know the semantics and typical uses of the four shuffle variants
- Implement warp-level stream compaction with ballot / popc
- Use the clearer abstraction cooperative groups provide
- Know the performance of atomics and how to avoid contention
The 32 threads in a warp execute in lockstep, and their registers physically live in the same register file. A shuffle instruction lets a thread read another thread's register in the same warp, with no trip through shared memory: no store, no barrier, no bank conflict, one instruction. This is the fastest thread-to-thread communication on a GPU.
| Primitive | Semantics | Typical use |
|---|---|---|
__shfl_sync(mask, v, srcLane) | Every thread reads srcLane's value | Broadcast |
__shfl_up_sync(mask, v, d) | Read lane - d | Prefix sum (scan) |
__shfl_down_sync(mask, v, d) | Read lane + d | Reduction |
__shfl_xor_sync(mask, v, m) | Read lane ⊕ m | Butterfly reduction, all-lane broadcast |
__ballot_sync(mask, pred) | Returns a 32-bit mask of who has pred true | Vote, stream compaction |
__activemask() | Returns the mask of currently active threads | Sync inside a divergent path |
Two warp reductions: down and xor
CUDA C++
1__inline__ __device__ float warpReduceSum(float v) {2 for (int off = 16; off > 0; off >>= 1)3 v += __shfl_down_sync(0xffffffff, v, off);4 return v; // only lane 0 holds the full sum5}6// off=16: lane0 += lane16, lane1 += lane17, ...7// off=8 : lane0 += lane8, lane1 += lane9, ...8// off=1 : lane0 += lane1 → doneCUDA C++
1__inline__ __device__ float warpAllReduceSum(float v) {2 for (int m = 16; m > 0; m >>= 1)3 v += __shfl_xor_sync(0xffffffff, v, m);4 return v; // all 32 lanes hold the full sum5}6// when every later thread needs the total (softmax normalization,7// for example), this version skips a broadcast.ballot + popc: stream compaction inside a warp
CUDA C++
1__global__ void compact(const int* in, int* out, int* count, int n) {2 int i = blockIdx.x * blockDim.x + threadIdx.x;3 int lane = threadIdx.x % 32;4 5 bool keep = (i < n) && (in[i] > 0);6 7 // one instruction: a bitmap of who in this warp wants to keep8 unsigned mask = __ballot_sync(0xffffffff, keep);9 10 // how many keepers sit before me? popc counts set bits11 int rank = __popc(mask & ((1u << lane) - 1));12 13 int base = 0;14 if (lane == 0) base = atomicAdd(count, __popc(mask)); // one atomic for the whole warp15 base = __shfl_sync(0xffffffff, base, 0); // broadcast to the warp16 17 if (keep) out[base + rank] = in[i];18}The trick in this snippet: 32 threads issue one atomic. If every thread called atomicAdd(count, 1) on its own, that address would serialize 32 ways. Ballot first to get the count and each thread's offset inside the warp, then one thread contends for the whole warp. Contention drops by 32x. This warp-aggregated atomics pattern shows up constantly in sparse work.
Cooperative groups: a clearer abstraction
CUDA C++
1#include <cooperative_groups.h>2#include <cooperative_groups/reduce.h>3namespace cg = cooperative_groups;4 5__global__ void reduceCG(const float* in, float* out, int n) {6 cg::thread_block block = cg::this_thread_block();7 cg::thread_block_tile<32> warp = cg::tiled_partition<32>(block);8 9 float sum = 0.0f;10 for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < n;11 i += blockDim.x * gridDim.x) sum += in[i];12 13 // the group owns the mask; you do not write 0xffffffff by hand14 sum = cg::reduce(warp, sum, cg::plus<float>());15 16 __shared__ float warpSums[32];17 if (warp.thread_rank() == 0) warpSums[warp.meta_group_rank()] = sum;18 block.sync(); // same as __syncthreads()19 20 if (warp.meta_group_rank() == 0) {21 sum = (warp.thread_rank() < warp.meta_group_size())22 ? warpSums[warp.thread_rank()] : 0.0f;23 sum = cg::reduce(warp, sum, cg::plus<float>());24 if (block.thread_rank() == 0) atomicAdd(out, sum);25 }26}- ▸Mask handled for you: no handwritten
0xffffffff, and you will not forget__activemask(). - ▸Composable granularity:
tiled_partition<8>cuts an 8-thread subgroup, more flexible than raw shuffle. - ▸Clearer semantics:
warp.thread_rank()says more thanthreadIdx.x % 32. - ▸Cost: some features need
-rdc=true, and grid-widegrid_group::sync()must be launched withcudaLaunchCooperativeKernel.
Atomics: skip them when you can
| Situation | Cost | What to do instead |
|---|---|---|
| Every thread in the grid atomic-adds the same address | Catastrophic, fully serial | Reduce inside the block first, one atomic per block |
| Every thread in a warp atomic-adds the same address | 32-way serial | Warp-aggregated atomics (ballot + one atomicAdd) |
| Atomics to scattered addresses (histogram) | Acceptable, completed in L2 | Hot buckets can accumulate in shared memory first |
atomicCAS spinlock | Deadlocks easily (when the warp diverges) | Switch to a lock-free algorithm or a reduction |
CUDA C++
1__global__ void histogram(const unsigned char* data, int* hist, int n) {2 // each block keeps a private histogram in shared memory, contention shrinks from the whole grid to the block3 __shared__ int local[256];4 5 for (int i = threadIdx.x; i < 256; i += blockDim.x) local[i] = 0;6 __syncthreads();7 8 for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < n;9 i += blockDim.x * gridDim.x) {10 atomicAdd(&local[data[i]], 1); // shared-memory atomic, an order of magnitude faster than global11 }12 __syncthreads();13 14 // then each block issues 256 global atomics15 for (int i = threadIdx.x; i < 256; i += blockDim.x) {16 if (local[i] > 0) atomicAdd(&hist[i], local[i]);17 }18}Self-check
Answer in your head first, then open the solution. If you cannot, reread that section.