本文目标

读完本文，你将能够：

  • 解析 GPU 线程束内的 Warp Divergence 物理成因及其解决方案
  • 定量分析 Shared Memory 版本与全局 Memory 收敛版本在小数据量下的假象等同（L1 Cache 效应）
  • 理解并实现 Thread Coarsening (线程粗化)，有效分摊调度税并打爆 HBM 内存带宽
  • 认知 FMA (Fused Multiply-Add) 指令融合与利用超大 L2 Cache (72MB) 实现的超物理理论峰值带宽机制

配套仓库：psmarter/CUDA-Practice
本系列文章与仓库中的源码、实验结果、目录结构相互对应；建议阅读文章时同步对照仓库中的实现。

对应代码路径

硬件环境：NVIDIA RTX 4090 (Ada Lovelace, sm_89)
128 SMs | FP32 82.6 TFLOPS | HBM 1008 GB/s | L2 72 MB | Roofline 拐点 81.9 FLOP/Byte

源文件 Kernel 名称 核心技术 测试规模
02_Reduction/01_reduce_sum/reduce_sum.cu simple_reduce_sum
convergent_reduce_sum
shared_reduce_sum		 发散消除 / 收敛索引 / Shared Memory 树状归约 N=2048
02_Reduction/02_reduce_optimized/reduce_optimized.cu segmented_reduce_sum
coarsened_reduce_sum
coarsened_reduce_max		 多 Block + atomicAdd、线程粗化 (COARSE_FACTOR=4)、Shared Memory 收尾 N=1048576 (1M)
02_Reduction/03_dot_product/dot_product.cu shared_dot_product
coarsened_dot_product
fma_dot_product		 点积 = 乘后归约、FMA 融合、L2 缓存热数据 N=1048576 (1M)

Kernel 名称与源码中 __global__ 函数签名完全一致。本篇多 Block 归约以 Shared Memory 树状折叠 + atomicAdd 收尾，未使用 __shfl_*；Warp 级无锁归约见 06 线程束原语与寄存器通信。

本篇在系列中的位置：承接 01 基础概念与分块 的带宽墙与 Shared Memory 直觉，将「分块与片上缓存」用于归约这一经典模式（多 Block、atomicAdd、线程粗化）。后续 03 前缀和与多块扫描 同属树形结构但需保留前缀和；06 线程束原语与寄存器通信 用 __shfl_* 做无 Shared Memory 的 Warp 归约；05 大模型算子与注意力归一化 的 Softmax/LayerNorm 依赖归约作为子步骤。

Baseline

问题陈述：将含有 NN 个元素的数组折叠为一个标量。这是深度学习 Softmax 或者 LayerNorm 前置的极限抽象。由于加法操作算术强度趋向于 0，该算子极度 Memory Bound。

Baseline 类别 测试场景 指标 值 数据来源
CPU 参考推演 (14核) N=1M (4.00 MB) Reduce Sum 耗时 4.69 ms [实测] Results/02_Reduction.md
CPU 参考推演 (14核) N=1M 点积 (8.00 MB) Dot Product 耗时 1.69 ms [实测] Results/02_Reduction.md
Naive Simple Reduce N=2048 极小规模 Divergence 耗时 0.0051 ms [实测] Results/02_Reduction.md
Segmented Naive GPU N=1M 大规模 多 Block 原子锁耗时 0.0084 ms [实测] Results/02_Reduction.md

瓶颈分析

为何朴素树状归约与分段归约无法打满物理带宽？原因解构如下：

  1. Warp Divergence 导致物理核心闲置浪费
    • Naive Reduce 中使用 stride = 1, 2, 4 配对。if (tid % stride == 0) 导致同 Warp 内活跃线程间距拉开。在 stride=16stride=16 时，单个指令提取单元下只剩 2 人执行，其余 30 个线程强行掩码阻塞，算力真空率高达 93.75% [理论]。
  2. 多 Block Segmented 同步隔离极高路费
    • 面对 1M 数组，按常规切分需激发 1024 个 Block（设极小规模双吃为 512）。1024 股数据在各自算出归约极值后，为了最终全图合并不受相互篡改，全数去冲撞排队 atomicAdd，自旋死锁造成严重阻塞。
  3. 低指令并行度 (ILP) 与发射税
    • 当单个线程只提取 1 或 2 个元素即发生一次 __syncthreads() 卡位时，底层指令流水由于没有连续无依赖的运算流填充，无法覆盖显存延迟本身。并且唤醒上千个短命 Block 占比过高。

优化思路

优化 1：Convergent Memory Indexing 解发散

解决的瓶颈：Warp 内部分步长断裂。
核心思想：彻底倒转树形折叠路线。stride 从 blockDim.x 开始折半倒退至 1。此时所有存活计算的 threadIdx.x 会严丝合缝地在 0…(N2−1)0 \dots (\frac{N}{2}-1) 之间完全连续。
预期收益：未活跃的 Warp 直接被调度器挂起免除执行，存活的 Warp 利用率恢复到 100%。

优化 2：Thread Coarsening 线程粗化

解决的瓶颈：极高频次的核间同步锁与低效 Kernel 调度比例。
核心思想：直接把原本指派给 1024 个短命 Block 做的事情，砍给仅有不到原来的四分之一或八分之一的粗粒度 Block 来做（如设定 COARSE_FACTOR=4）。每个线程不加掩饰地连续吞下并内部私有化 sum += input[i] 多达 8 此，将其直接拦截在最内层无等待寄存器端（提升 ILP）。
预期收益：极大压缩 Block 并发池总量，将 1M 归约全带宽吞没打入 0.0047 ms 内 [实测]。

优化 3：FMA 底层融合与 L2 跨越

解决的瓶颈：点积中乘法与加法的双周期分裂损耗。
核心思想：使用 fmaf(a, b, sum) 将独立的相乘、相加合二为一，共用一级硬件流水以节省周期并增加舍入精度。外加上，如果我们故意使双阵列容量（2×4MB=8MB2 \times 4\text{MB} = 8\text{MB}）控制在极高配置的 72MB RTX 4090 L2 Cache 之内进行热重载，就能打碎物理位阶，触发直对 SM 的爆表读取。
预期收益：测算出破表理论宽度的 1506 GB/s，完全压倒 HBM 发车极限 [实测]。

关键代码解释

Divergence 修复的几何折转

1
2
3
4
5
6
7
8
9
10
// 来源：02_Reduction/01_reduce_sum/reduce_sum.cu（convergent_reduce_sum / shared_reduce_sum 的 stride 逻辑）
    // 【错误模式】simple_reduce_sum：stride 从 1 倍增，if (threadIdx.x % stride == 0) 导致 Warp 内活跃线程不连续

    // 【收敛模式】：stride 从 blockDim.x 减半至 1，活跃线程 tid 始终落在 [0, stride)，即连续无间隙
    for (int stride = blockDim.x / 2; stride >= 1; stride /= 2) {
        __syncthreads();  // shared 版本先同步再读
        if (threadIdx.x < stride) {
            shared_data[threadIdx.x] += shared_data[threadIdx.x + stride];
        }
    }

要点解读：

  • 虽然同为 log⁡N\log N 次加法，收敛版本在 stride≤16stride \le 16 即倒数第四轮步长衰减中，由于全集只剩不到 16 个线程在工作，它会全数合并落在唯一的 Warp 0 手中，其余 Warp 早已安全撤退退出轮询池。

线程粗化的 Register 化截留

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
// 来源：02_Reduction/02_reduce_optimized/reduce_optimized.cu : coarsened_reduce_sum 核心片段
__global__ void coarsened_reduce_sum(float* input, float* output, int length) {
    __shared__ float shared_data[BLOCK_SIZE];
    int tid = threadIdx.x;
    int sid = 2 * COARSE_FACTOR * blockDim.x * blockIdx.x + tid;

    float sum = 0.0f;
    if (sid < length) {
        sum = input[sid];
        for (int i = 1; i < COARSE_FACTOR * 2; ++i) {
            if (sid + i * BLOCK_SIZE < length)
                sum += input[sid + i * BLOCK_SIZE];
        }
    }
    shared_data[tid] = sum;
    // 树状 Shared Memory 归约 + 最后 atomicAdd(output, shared_data[0])
    for (int stride = blockDim.x / 2; stride > 0; stride /= 2) {
        __syncthreads();
        if (tid < stride) shared_data[tid] += shared_data[tid + stride];
    }
    if (tid == 0) atomicAdd(output, shared_data[0]);
}

要点解读：

  • COARSE_FACTOR * 2（=8）将每线程负责的元素数拉大到 8 个，先在寄存器 sum 中串行累加，再写入 shared_data[tid]，Block 数降为约 1/8，__syncthreads() 与 atomicAdd 调用次数同步减少。收尾仍为 Shared Memory 树状归约 + 单次 atomicAdd，本实现未使用 __shfl_*。

结果与边界

性能对比

测试条件：RTX 4090 (sm_89sm\_89) , 参数 100 迭代求均
数据来源：Results/02_Reduction.md 原始实机日志

1. 极小数据集 N=2048 的无力感

并行实现手段 执行时间 L1/L2 效应评价 数据性质
Naive Simple Reduce 0.0051 ms 单纯依赖 Global 抖动 [实测]
Convergent Reduce 0.0038 ms 完美贴合 L1 局部热区 [实测]
Shared Memory Reduce 0.0038 ms 未拉开与前者的实质差别 [实测]

在这个规模仅占据 8KB8\text{KB} 容量的数据上，把数据显性写进 Shared Memory 的动作与由于全连续访存被 128KB L1 Cache 全面隐性截获的 Convergent Global 版速度一模一样。底层硬件的 L1 代打抵消了人力干预。

2. 宏观数据集 N=1,048,576 (1M) 粗化降维战

测试环境 Total Kernel 耗时 对比基数 带宽折现 数据性质
CPU 参考 (14核) 4.69 ms 1.00x - [实测]
GPU Segmented (多包细切) 0.0084 ms 558.33x 476.19 GB/s [实测]
GPU Coarsened (寄存器粗卷) 0.0047 ms 991.52x 887.48 GB/s [实测]

当将 COARSE_FACTOR=4 外挂至核心体系内，通过在最内核强吃 8 倍元素的方法直接缩除掉了 87.5% 的跨 Block 落锁频次与 Kernel 线程启爆数量！887.48 GB/s 的实跑总线已经摸平到 4090 极值带宽的 88%。对于 FLOP/Byte=0.125 的无计算量操作，这已是该卡在此问题物理尺寸上的顶峰。

边界条件与局限

  • L2 穿障极限 (The Fake 1506 GB/s)：在进行 Dot Product 实验中测出的 0.0056 ms -> 1506.49 GB/s 总量，远超 4090 总线极值。这种现象只会出没于测试体量（8 MB）完全坍缩于其标定超大的 72MB L2 范畴内。当 N 飙升至千万击穿 L2 Threshold 后，所有的魔法都会打回原形回落。
  • 系数不能无脑拉大：COARSE_FACTOR 若激增至 32 以上，将发生灾难性的 Register Spilling 溢出至 Local Memory 反噬耗时。

常见误区

  1. 误区：在当前高架构（如 sm_89）下，纯 CPU 代码加上 -O3 -mavx2 就可以轻松与微小型 GPU Kernel 抗衡。
    实际：在所有我们构建出的测试用例中，哪怕是只处理极细的一兆数组，只要你动用了粗化将吞地拉展，GPU 依然能爆出超越主流高核心桌搭级 CPU 千倍 (991+倍) 的绝对降维制裁表现 [实测]。
  2. 误区：代码里强制用 fmaf(a, b, c) 一定比 a*b + c 跑得更带感。
    实际：在最高优化的 NVCC 管道中毫无分别 [实测 0.0056ms对锁 0.0056ms]。当代现代编译器面对简单线性多项式，早就自行替程序员做出了 FMA 底层硬路由融合，它更多是规避特殊舍入的修饰语。

系列导航

前置阅读

文章 与本篇的衔接
01 基础概念与分块 建立带宽墙、Roofline 与 Shared Memory Tiling 直觉；本篇在同一存储层级上做「归约」模式并引入多 Block 与 atomicAdd

推荐后续（承上启下）

文章 与本篇的衔接
03 前缀和与多块扫描 同属树形/线性扫描结构，但需保留中间前缀结果，多 Block 与分段策略与归约形成对照
06 线程束原语与寄存器通信 本篇用 Shared Memory + atomicAdd 收尾；06 用 __shfl_* 在 Warp 内无 Shared Memory 完成归约/Scan，进一步压延迟
05 大模型算子与注意力归一化 Softmax、LayerNorm、RMSNorm 均以归约（max/sum/方差）为子步骤，本篇为理解其 Kernel 打底

顺序导航