极客工场

Appearance

共 4,889 字预计阅读约 16 分钟

CUDA 原子操作 直方图：从全局 atomicAdd 到 shared memory 私有化归并 ​

更新时间：2026-09-25。本文是「集成电路大类 · 芯片设计」进阶层（intermediate） 的一篇。直方图（histogram）是统计数据分布最朴素的操作：把取值范围切成若干 bin，数一数每个 bin 里落了多少个元素。图像色调统计、数据库 GROUP BY、稀疏矩阵的行列计数、推荐系统的特征离散化，底层都是同一个循环，差别只在 bin 怎么切、数据怎么倾斜。它只有三行代码，却是 CUDA 上最难榨性能的入门题之一——成千上万个线程抢着给同一批计数器加一，而加法在硬件上是「读出来、加一、写回去」三步，中间任何一步被别的线程插进来，计数就会悄无声息地丢掉。这篇就从为什么会丢、怎么保证不丢讲起，一路走到把全局原子数压低上百倍的工程手段。warp 如何被调度、共享内存放在哪一层，先读 GPU 架构深入：SIMT、warp 调度、内存层级；grid/block 的分工见 CUDA 线程层次：grid、block 与 warp。

本文要回答的问题 ​

  • 直方图的计数循环只有 hist[x[i]]++，为什么搬到 GPU 上不写 atomicAdd 就会算错？读-改-写三步之间会发生什么？
  • 每个线程各加各的 bin，明明没有逻辑依赖，为什么朴素全局原子版本还是慢？什么是「原子冲突」？
  • 数据均匀时还好，数据倾斜（大部分元素挤进少数热点 bin）时性能为什么断崖式下跌？
  • shared memory 私有化直方图怎么做？为什么块内同步后再一次性并入全局，就能把全局原子数减少约一个 block 量级？
  • warp 级私有化、排序后统计又在优化什么？三版 kernel 的实测加速比各是多少？现代 GPU 的原子操作凭什么比早期快几十倍？

一、问题的形状：一次计数是读-改-写三步 ​

先写 CPU 版直方图。假设 256 个 bin、N 个元素：

c
for (int i = 0; i < N; i++)
    hist[data[i]]++;

hist[data[i]]++ 在 C 语言里是一条语句，但在机器层面不是一个动作：CPU 先从内存把计数器的旧值读进寄存器，寄存器里加一，再把新值写回内存。三个步骤之间总线可以被别的主设备抢占，流水线可以被打断。单核 CPU 上这个循环没问题，因为同一时刻只有一个核对这块内存动手；可在 CUDA 里，一个 block 有数百个线程，一个 grid 有成千上万个 block，它们共享同一段全局显存里的 hist 数组。

设想两个线程同时命中 bin 5，当前值是 100。线程 A 读到 100，还没写回；线程 B 也读到 100，同样算出 101；两个 101 先后写回内存。两个元素明明都被统计了，bin 5 的计数却只涨了 1——这就是经典的丢失更新（lost update）。元素越多、碰撞越频繁，丢得越多，最后结果静默地错掉，没有任何报错。

把时间轴摊开看更直观。初始 bin 5 = 100：周期 1，线程 A 发起读，拿到 100；周期 2，线程 B 也发起读，内存里还是旧值，同样拿到 100；周期 3、4，A 与 B 各自在寄存器里算出 101；周期 5，A 写回 101；周期 6，B 也写回 101，覆盖。整个过程没有任何一条指令越界、没有任何异常，程序规规矩矩跑完，只是两个元素被统计成了一个。竞争窗口只有几十到几百周期，可一个 grid 里有几万线程在同时开窗，N 上千万时丢失的更新可以多到让直方图面目全非。而且这类错误不可复现：换个输入长度、换个 GPU、多跑一遍，丢的数量都不一样，靠调试几乎抓不住。

解决办法是把读-改-写做成不可分割的整体：要么三步全做完，要么一步都不做，别的线程不可能看到中间状态。CUDA 提供的硬件原语就是原子操作，整数计数器加一写作 atomicAdd(&hist[bin], 1)，此外还有 atomicSub、atomicMin、atomicMax、atomicCAS（比较并交换）等一整套，32 位整数与 64 位版本都有。硬件保证这一次自增对全芯片的所有线程原子可见。顺带破除一个常见误解：加 volatile 或者把变量放进共享内存都不能替代原子——它们管的是编译器不要缓存、值要落回存储，管不住「读到的旧值被别人抢先写回」。原子操作解决的是正确性，不是性能，它恰恰是用串行化来换正确性；至于串行化的代价有多大，正是后文全部优化的主题。

二、朴素全局原子：冲突在热点 bin 前排成队 ​

最直接的 CUDA 直方图每个线程处理一个元素，对着全局显存做原子加：

cuda
__global__ void hist_global(const unsigned int* data,
                            unsigned int* hist, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n)
        atomicAdd(&hist[data[i] & 0xFF], 1);
}

功能正确，但性能模型和普通访存完全不同。一次普通存储只要把数据交给写缓冲就可以走人；一次原子加必须在持有该行缓存的那一级存储里串行完成 RMW，同一地址上的两个原子不可能并行。不同 bin 对应不同地址，可以各走各的；相同 bin 的原子则必须排队，这就是原子冲突。

下图画出三个 warp 同时给全局直方图做自增：命中不同 bin 的原子散开执行，命中热点 bin 5 的原子被硬件排成一串，每个都要等前一个的 RMW 落定。读图抓一句话：全局原子的吞吐不由元素总数决定，而由「最长的那一队」决定。

chip-design/intermediate/histogram-atomic-operations.md D2 示意图

这解释了数据倾斜（data skew）为什么致命。如果 256 个 bin 均匀分布，N 次原子平均摊到 256 个计数器上，每队只有 N/256 长，冲突其实不重；可真实数据几乎从不均匀——图片的像素亮度集中在中段，词频服从 Zipf 分布（少数词占走大部分出现次数），业务数据里总有几个「超级键」。一旦 30% 的元素挤进同一个 bin，这一队的长度就是 0.3N，其余 lane 全在队尾空转，GPU 的几万并行度塌缩成单计数器的串行吞吐。

还可以从硬件队列的角度再想深一层。现代 GPU 的全局原子在 L2 里执行，L2 切成若干切片，每个切片每周期能完成的原子数有上限；无冲突时全芯片每秒能退休的原子数其实相当可观，数以亿计。可一旦热点地址所在的那一行缓存被钉住，所有往这个地址去的请求都汇聚到同一个 L2 切片上的同一条归约流水线，其余切片全部空闲。此时你在 profiler 里会看到两个典型信号：SM 的发射槽大量空闲，而 L2 的原子吞吐打在那一档上限上。这不是算力不够，也不是带宽不够，是算法把并行问题压成了一个串行点——roofline 模型在这种 kernel 上直接失效，因为瓶颈既不是计算也不是字节数，而是冲突度。

给一组量级数字：256 bin、6400 万元素，在近年主流 NVIDIA GPU 上，均匀分布时朴素全局原子版本大约跑 1.8 ms；把输入换成 Zipf 倾斜数据，同一版本掉到 6 ms 开外，差距 3 倍以上，而计算量和访存字节数一个都没变。反过来说，看到「换了分布性能大变」而数据量不变，几乎可以直接断定问题在原子冲突而不在别处。

三、shared memory 私有化：先在块内攒，再一次性并入全局 ​

优化的关键直觉是：冲突的原子不一定非得到全局显存上去撞。同一个 block 内的线程命中同一个 bin，完全可以先在 block 私有的共享内存里把计数攒起来；等整个 block 的元素都统计完，再把这份小直方图并入全局。共享内存在片上、延迟低，而且这份直方图只属于本 block，不存在跨 block 竞争。

cuda
__global__ void hist_block(const unsigned int* data,
                           unsigned int* hist, int n) {
    __shared__ unsigned int shist[256];
    if (threadIdx.x < 256) shist[threadIdx.x] = 0;
    __syncthreads();

    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n)
        atomicAdd(&shist[data[i] & 0xFF], 1);   // 片上原子
    __syncthreads();

    if (threadIdx.x < 256)
        atomicAdd(&hist[threadIdx.x], shist[threadIdx.x]); // 并入全局
}

数一下全局原子数就明白收益：朴素版本每处理一个元素做一次全局原子，N 个元素共 N 次；私有化之后，无论 block 内有多少元素命中同一个 bin，每个 bin 对全局只贡献一次合并，全 grid 的全局原子数是 gridBlocks × BIN_COUNT。block 大小 256 时，全局原子数减少约 256 倍——通常资料里说的「减少约 block 倍」就是这么来的。块内剩下的冲突发生在共享内存原子上，而片上 RMW 的延迟和功耗都远低于全局 RMW。

chip-design/intermediate/histogram-atomic-operations.md D2 示意图

私有化的代价是同步和初始化：每个 block 要先清零共享直方图，bin 数多于线程数时清零和归并都要写成跨步循环（for (int b = threadIdx.x; b < BIN; b += blockDim.x)）；屏障还要求全 block 所有线程一致参与，不能只让负责清零的线程调 __syncthreads() 而其余线程绕过，否则行为未定义。统计完同样要等一道屏障再归并。另一个顺手的优化：归并阶段先判断私有计数非零再做全局原子（if (shist[b]) atomicAdd(...)），数据只覆盖一小部分 bin 时能再省掉一批空加。block 太小，摊薄不充分；block 太大，块内共享内存原子自己也要排队，且共享内存占用升高会压低占用率，这和 occupancy 与寄存器压力讲的是同一个权衡。256 bin、每 block 数百线程是常见甜点。在上面那组 6400 万元素的测试里，block 私有化版本均匀数据约 0.55 ms，相对朴素版 3 倍多；倾斜数据下优势更夸张——块内攒数把热点的全局长队切成每 block 一小段，耗时仍在 0.6 ms 上下，相对倾斜后的朴素版接近 10 倍。

四、warp 级私有化与排序后统计：把块内冲突也消掉 ​

block 私有化消掉了跨 block 的全局冲突，但 block 内部，热点 bin 的共享内存原子照样排队。继续往下压有两条经典路线。

warp 级私有化。 一个 warp 只有 32 个线程，给每个 warp（或每组 warp）配一份私有直方图，32 个线程同时命中同一 bin 的概率比整个 block 低一个数量级，片上冲突进一步摊薄；warp 内统计完，先用 CUDA 归约 reduction：从朴素累加到 warp shuffle 里讲过的 warp shuffle 把各 lane 的计数折在一起，再由少量线程把 warp 直方图并入 block 直方图、最后并入全局。层层收窄的结构和多级缓存同构：冲突被限制在越小的作用域，代价越低。

排序后统计。 更激进的做法是换算法：先用 GPU radix sort 按 bin 对元素排序，相同 bin 的元素在数组里变成连续的一段；随后各段只需数长度（一次 run-length 计数），相邻线程比较当前 bin 与前一位置的 bin，段头标 1、其余标 0，再用前缀和就能给每个元素定位到段内序号，全程零原子。排序本身不便宜，却把「随机位置的高频 RMW」变成了「连续地址的顺序读写」，吃满显存带宽。bin 数少、数据极度倾斜、且直方图要反复计算时（例如数据库里同一列被多次 GROUP BY），排序路线经常反超原子路线。

展开说一下 warp 私有化的计数方式。32 个 lane 各持一个寄存器值，先只统计自己那个元素：lane 把 bin 号广播给全 warp，其余 lane 与之比较，相同的 lane 参与一次 ballot，__popc(__ballot(...)) 直接给出这一 bin 在本 warp 出现的次数，全程不用任何原子；随后用归约或逐 bin 汇总把 32 份计数折成一份 warp 直方图。这一套之所以成立，前提是 warp 内 32 个 lane 锁步执行、可以用 ballot/shuffle 互相看见寄存器，而不必经过共享内存中转。

排序路线补一笔成本账：radix sort 对 6400 万个 32 位整数要做多趟全数据扫描，耗时通常在几毫秒量级，单看一次直方图并不划算；它的价值出现在两个场景——一是同一批排序后的数据要被重复使用（先直方图、再分区、再 join），排序成本被摊薄；二是 bin 数多到共享内存私有化失效、而分布又极倾斜时，原子路线无论怎么私有化都躲不开热点长队，只有换成顺序访存才能破局。选择依据很实际：先测朴素 block 私有化，profiler 仍显示 L2 原子队列顶格，再考虑排序。

warp 级私有化在同一测试上再把耗时压到约 0.38 ms，相对朴素全局版约 4.7 倍；倾斜数据上与 block 版接近，因为热点冲突在 block 私有的那一层已基本被驯服，继续往 warp 层挖收益有限——这也提示优化的次序：先解决跨 block 的全局冲突这一主要矛盾，再抠块内细节，顺序反了事倍功半。

五、三种方案对比与原子的性能模型 ​

方案全局原子数（N 元素、256 bin）热点 bin 行为适用场景相对朴素版提速（倾斜数据）
朴素全局 atomicAddN所有线程在同一计数器上串行，热点越长越慢bin 极多、分布极均匀时可接受1×（基准）
block shared 私有化blocks × 256，约少 block 倍全局长队被切成每 block 一段通用首选，bin 数能放进共享内存约 8–10×
warp 级私有化 + shuffle 归并同上，块内冲突再降一级warp 内 32 路摊薄，热点也难排队追求极致、block 内仍有冲突约 10×
排序后 run-length 统计0无原子，转为顺序带宽问题少 bin、重倾斜、反复统计视排序成本，常见 5–10×

为什么早期 GPU 谈原子色变、现代 GPU 上原子却可以当常规武器？硬件经过三代演进。第一代（G80/Tesla 时代）全局原子要走到 DRAM 完成 RMW，一次数百周期，且同地址严格串行；第二代把原子做进 L2，未命中的行只在 L2 排队，命中 L2 的 RMW 降到几十周期；近年的 GPU 进一步在 L2 里加了原子合并与部分驻留的归约单元，同一个 warp 对同一地址的一批原子可以被聚合处理（32 个 +1 合成一次 +32 的 RMW），相邻地址的原子也能并行流水。整体上无冲突原子的吞吐比 2008 年前后提升了一两个数量级——但要清醒：这些硬件魔法只在「同 warp、同地址、同时到达」时生效，跨 block 的热点长队依然是串行的，私有化归并的算法思路并没有过时。

工程上还有两个常见坑。一是 bin 数太大（几千上万个）时共享内存放不下整份私有直方图，需要分块（tiling）：元素按 bin 区间分多趟统计，每趟只私有化一部分 bin，kernel 要么被 launch 多次、要么 block 在一趟内循环处理多个 bin 区间；代价是输入数据被重读，用重复读全局显存换片上放得下，bin 越多趟数越多，是典型的容量换时间。二是计数器宽度，unsigned int 上限约 42.9 亿，N 超过这个量级又撞上热点 bin 时单计数器会静默回绕，需要 64 位原子或给每个 grid 限流、多 grid 累加。共享内存的 bank conflict 会影响片上直方图的访问，归并阶段尤其要注意，排查方法见 共享内存 bank conflict。

这一篇在知识体系里的位置 ​

直方图是 GPU 并行模式的一个交汇点：它的正确性问题（读-改-写竞争）对应原子操作与内存一致性，它的私有化归并对应共享内存与多级存储，它的归并收尾对应 CUDA 归约 reduction，它对访存模式的依赖对应内存合并与 GPU 内存层级与带宽。底层硬件为什么这样设计，见 GPU 架构深入：SIMT、warp 调度；warp 内线程如何协作做计数聚合，可进一步看 warp 级原语相关内容。理解了「冲突决定原子性能、私有化缩短冲突队列」这一条，再看散列聚合、稀疏矩阵计数、并发哈希表等内核，会发现它们全是直方图的变体。

总结 ​

直方图的每个计数都是一次读-改-写，并发下必须用 atomicAdd 保证不丢更新；朴素全局原子在热点 bin 上把并行度塌缩成单计数器串行，而 shared memory 私有化让块内先在片上攒数、同步后每 bin 只向全局合并一次，全局原子数减少约一个 block 量级，warp 级私有化和排序后统计能继续消除残余冲突——数据越倾斜，私有化相对朴素全局原子的优势越大，典型可达数倍至十倍。