<<<CUDA C++ 学习路线grid · block · warp · lane
EN中文
阶段 4 · 深入底层11 / 15约 24 分钟cuda/09_warp_shuffle.cu

Warp 原语：绕过共享内存的数据交换

shuffle、ballot、cooperative groups，以及原子操作的正确用法

学完这一课你会
  • 掌握四个 shuffle 变体的语义和典型用法
  • 用 ballot / popc 实现 warp 内的流压缩
  • 了解 cooperative groups 提供的更清晰抽象
  • 知道原子操作的性能特征和规避争用的技巧

warp 内的 32 个线程锁步执行，它们的寄存器物理上就在同一个寄存器文件里。shuffle 指令让线程可以直接读取同一 warp 内其他线程的寄存器，完全不经过共享内存——没有存储、没有屏障、没有 bank conflict，一条指令搞定。这是 GPU 上最快的线程间通信方式。

原语语义典型场景
__shfl_sync(mask, v, srcLane)所有线程读取 srcLane 的值广播
__shfl_up_sync(mask, v, d)读取 lane − d 的值前缀和（scan）
__shfl_down_sync(mask, v, d)读取 lane + d 的值归约
__shfl_xor_sync(mask, v, m)读取 lane ⊕ m 的值蝶形归约、全 lane 广播
__ballot_sync(mask, pred)返回 32 位掩码，标记谁的 pred 为真投票、流压缩
__activemask()返回当前活跃线程的掩码分叉路径内的同步

两种 warp 归约：down 与 xor

shfl_down：结果只在 lane 0
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;      // 只有 lane 0 拿到完整的和5}6//  off=16: lane0 += lane16, lane1 += lane17, ...7//  off=8 : lane0 += lane8,  lane1 += lane9,  ...8//  off=1 : lane0 += lane1   → 完成
shfl_xor：所有 lane 都拿到结果（蝶形）
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;      // 32 个 lane 全都持有完整的和5}6// 后续每个线程都要用到总和时（比如 softmax 归一化），7// 这个版本省掉了一次广播。

ballot + popc：warp 内流压缩

把满足条件的元素紧凑地写出去
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    // 一条指令拿到「本 warp 里谁要保留」的完整位图8    unsigned mask = __ballot_sync(0xffffffff, keep);9 10    // 我前面有几个人要保留？popc 数二进制 1 的个数11    int rank = __popc(mask & ((1u << lane) - 1));12 13    int base = 0;14    if (lane == 0) base = atomicAdd(count, __popc(mask));   // 整个 warp 只做 1 次原子15    base = __shfl_sync(0xffffffff, base, 0);                // 广播给全 warp16 17    if (keep) out[base + rank] = in[i];18}

这段代码的精妙之处在于：32 个线程只发起了一次原子操作。如果每个线程各自 atomicAdd(count, 1)，同一地址上会有 32 路串行争用。先在 warp 内用 ballot 算出总数和各自的偏移，再由一个线程代表整个 warp 去抢，争用降低 32 倍。这个「warp 聚合原子（warp-aggregated atomics）」模式在稀疏计算里极其常用。

Cooperative Groups：更清晰的抽象

用 cg 重写归约
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    // 掩码由 group 自己管理，不用手写 0xffffffff14    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();                                   // 等价于 __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}
  • ▸掩码自动管理：不用再手写 0xffffffff，也不会忘记 __activemask()。
  • ▸粒度可组合：tiled_partition<8> 能切出 8 线程的子组，比裸 shuffle 灵活。
  • ▸语义清晰：warp.thread_rank() 比 threadIdx.x % 32 表意明确得多。
  • ▸代价：需要 -rdc=true 编译（部分特性），且 grid 级同步 grid_group::sync() 要用 cudaLaunchCooperativeKernel 启动。

原子操作：能不用就不用

场景代价对策
全 grid 线程原子加同一地址灾难级，完全串行先做 block 内归约，每 block 只原子一次
warp 内原子加同一地址32 路串行warp 聚合原子（ballot + 单次 atomicAdd）
原子加分散地址（如直方图）尚可，L2 上完成热点桶可先在共享内存里累加
atomicCAS 自旋锁极易死锁（warp 内分叉时）改用无锁算法或归约
直方图：共享内存私有化
1__global__ void histogram(const unsigned char* data, int* hist, int n) {2    // 每个 block 先在共享内存里维护一份私有直方图，争用范围从全 grid 缩到 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);      // 共享内存原子，比全局快一个量级11    }12    __syncthreads();13 14    // 最后每个 block 只对全局做 256 次原子加15    for (int i = threadIdx.x; i < 256; i += blockDim.x) {16        if (local[i] > 0) atomicAdd(&hist[i], local[i]);17    }18}

自测

先自己在心里回答一遍，再展开对照。答不上来的说明这一段值得重读。