算子在做什么，先看一个算例
这一篇不是一个神经网络算子，而是一个底层原语：atomicAdd——多个线程往同一个地址做加法时，保证结果不丢的原子操作。
它出现在所有“多对一”的累加场合：reduce 的收尾、稀疏 embedding 的梯度累加、直方图统计、bias 聚合。这篇文章要问的是一个很具体的问题：当累加的量是 half（2 字节）时，为什么 atomicAdd 会变得很慢，以及怎么让它快起来。
先看 atomicAdd 为什么不能省
普通的 c[i] = c[i] + v 是“读—改—写”三步。如果两个线程同时做这件事，它们可能都读到同一个旧值、各自加完再写回，后写的那个会盖掉先写的——一次加法凭空消失。atomicAdd 把这三步做成一个不可分割的整体，由硬件在 L2 上完成：同一个地址的操作会排队串行，但每个操作一定生效。
代价也随之而来：同一地址的原子操作吞吐，取决于每次搬运的字节数。硬件通路按 32 位粒度组织，一次搬 2 字节就是只填了四分之一。这篇文章的三个版本，做的就是同一件事的三种做法：把每次原子操作的宽度凑到 4 字节。
测例：一个刻意造出来的最坏情况
测例是 half 向量内积 c = Σ a[i]*b[i]，但重点不在内积本身。数据规模 N = 32×1024×1024 = 33,554,432，启动配置是 131072 个 block × 256 线程——线程总数恰好等于 N。
一个必须知道的坑：结果是 32，不是 335544.32
这个测例的数学真值是 33,554,432 × 0.1 × 0.1 = 335,544.32。但三个版本打印出来的是 32、64、32。这不是 bug，是 half 的精度问题，而且它和原子加的宽度毫无关系——三个版本都有这个问题。
一个用 half 累加的内积，无论原子加优化得多快，算出来的都是错的值。所以下面性能一节的数字，衡量的是原子操作的吞吐，不是这个内积算得对不对。三个版本的对拍方式也因此只能是“对拍打印值是否与预期一致”，仓库里本来就没有正确性校验代码。
代码逐段解析
三个版本共享同一个 kernel 骨架——网格步长循环 + 每个线程最后发一条原子指令——差别只在那条原子指令怎么写。
① v0：基线，2 字节的 half 原子加
// 【函数功能】两阶段内积 c = Σ a[i]*b[i]：
// 阶段① 每线程把分给自己的 a[gid]*b[gid] 累进寄存器 temp；
// 阶段② 每线程把自己那份 temp 用一条 atomicAdd 汇入同一个地址 c。
__global__ void dot(half* a, half* b, half* c, int n) {
const int nStep = gridDim.x * blockDim.x; //【grid-stride 步长】nStep = 总线程数 = N
half temp = 0.0; // 【线程私有的部分和】从 half 常量 0.0 起步
int gid = blockIdx.x * blockDim.x + threadIdx.x;
while (gid < n) {
temp = temp + a[gid] * b[gid];
gid += nStep; // 跳到下一个归本线程的元素
}
atomicAdd(c, temp);// 【本版的主角】2 字节宽的原子加：33,554,432 个线程全往 c 这一个地址发 atom.add.noftz.f16
}第 27 行 atomicAdd(c, temp) 就是全部。c 是 half*，所以这一条会编译成 PTX 的 atom.add.noftz.f16——一次改 2 字节。
原生支持情况要分清楚（这两个门槛不一样，很容易记混）：
本机是 sm_120，两条路径都是硬件原生实现，不经过 CAS 模拟。atom.add.noftz.f16 里的 noftz 是“不清除非规格化数”——对 half 的加法语义有细微影响，但不是本文关心的重点。
CUDA 文档里有明确说明：对 __half2 的原子操作，原子性是对 两个 half 分量分别成立的，不保证整个 32 位作为一个整体原子。这一点在下一节的 v2 里会变得很关键——因为 v2 正是靠“给邻居 lane 加 0”来凑宽度的，它依赖于两个分量互不干扰。
② v1：改数据视图，把两个 half 打包
// 数组排布就可能落在奇数地址上，后面 reinterpret_cast 成 half2* 就破坏对齐。
template<typename T, size_t pack_size>
struct alignas(sizeof(T) * pack_size) Pack {
T elem[pack_size];
};
template<typename T, int32_t pack_size>
__device__ __inline__ void AtomicAdd(Pack<T, pack_size>* address,
T val) {
#pragma unroll
for (int i = 0; i < pack_size; ++i) {
atomicAdd(reinterpret_cast<T*>(address) + i, static_cast<T>(val));
}
}
template<>
__device__ __inline__ void AtomicAdd<half, 2>(Pack<half, 2>* address, half val) {
half2 h2_val;
h2_val.x = static_cast<half>(val);
h2_val.y = static_cast<half>(val);
atomicAdd(reinterpret_cast<half2*>(address), h2_val);
}思路很直接：既然一条原子指令只有 2 字节，那就让每次原子操作管两个元素。做法是换一套访存视图——不再把数组当 half*，而是当 Pack<half, 2>*（两个 half 一组），这样相邻两个 half 在同一个 4 字节槽里，一条 atomicAdd(__half2*) 就能一起改。
alignas(sizeof(T) * pack_size) 又是那个关键开关——它把 Pack<half, 2> 的对齐抬到 4 字节，否则 reinterpret_cast<half2*> 就可能在奇数地址上，破坏对齐。
__global__ void dot(Pack<T, pack_size>* a, Pack<T, pack_size>* b, Pack<T, pack_size>* c, int n){ // 实例化：T=half、pack_size=2（见 main 里的 dot<half, 2><<<...>>>）
const int nStep = gridDim.x * blockDim.x; // 线程总数 = N，组数 = N/2 → 多数线程空转一圈
T temp = 0.0;
int gid = blockIdx.x * blockDim.x + threadIdx.x;
while (gid < n / pack_size) { // 只遍历组；n 为奇数时会漏掉末尾元素（见文件头"长度对齐"）
for (int i = 0; i < pack_size; i++) { // 【一组 2 个元素都乘进来】a[gid].elem[i] 就是原来第 2*gid+i 个 half
temp = temp + a[gid].elem[i] * b[gid].elem[i];
}
gid += nStep;
}
AtomicAdd<T, pack_size>(c, temp); // 【每个线程一次 half2 原子加】后一半线程在这里加了个 0，次数与 v0 持平
}第 54 行的循环上界变成了 n / pack_size——只遍历“组”，每组两个元素都在循环体里乘进来（第 55-57 行）。
两个硬约束：数组长度必须能被 2 整除，而且基地址必须 4 字节对齐。这两个要求意味着 v1 不是一个能直接丢进通用库的方案——它要求调用方保证数据布局，而真实场景里的缓冲区往往是从别处切出来的一段，未必满足。这正是 v2 要解决的问题。
③ v2：不动数据，只在对齐上做文章
template<typename T, typename std::enable_if<std::is_same<half, T>::value>::type* = nullptr>
__device__ __forceinline__ void FastSpecializedAtomicAdd(T* base, size_t offset,
const size_t length, T value) {
#if ((defined(CUDA_VERSION) && (CUDA_VERSION < 10000)) \
|| (defined(__CUDA_ARCH__) && (__CUDA_ARCH__ < 700)))
atomicAdd(reinterpret_cast<half*>(base) + offset, static_cast<half>(value));
#else
// Accounts for the chance base falls on an odd 16 bit alignment (ie, not 32 bit aligned)
__half* target_addr = reinterpret_cast<__half*>(base + offset);
bool low_byte = (reinterpret_cast<std::uintptr_t>(target_addr) % sizeof(__half2) == 0);
if (low_byte && offset < (length - 1)) {
__half2 value2;
value2.x = value;
value2.y = __float2half_rz(0);
atomicAdd(reinterpret_cast<__half2*>(target_addr), value2);
} else if (!low_byte && offset > 0) {
__half2 value2;
value2.x = __float2half_rz(0);
value2.y = value;
atomicAdd(reinterpret_cast<__half2*>(target_addr - 1), value2);
} else {
atomicAdd(reinterpret_cast<__half*>(base) + offset, static_cast<__half>(value));
}
#endif这一版的思路完全不同：数据和视图都不动，只在原子加调用的那一刻判断“目标地址是不是 4 字节对齐”，能凑成就用 4 字节，凑不成退回 2 字节。
关键在第 1、2 分支的“给邻居加 0”：因为 __half2 的原子性对两个分量分别成立，给邻居 lane 加 0 不会改变邻居的值，所以这个操作既拿到了 4 字节的宽度，又没有副作用。分支 3 是兜底：当目标正好是首元素且地址不对齐、或者末元素且右边没格子时，凑不成对，就老老实实做一次 2 字节加法——正确性优先，快慢其次。
本测例的调用是 FastAdd(c, 0, 2, temp)：offset=0、length=2，而 cudaMalloc 返回的地址必然是 4 字节对齐的，所以恒定走分支 1。这也意味着 c[1] 全程被动地加了一路 0——白占了一半事务。
template<typename T, typename std::enable_if<!std::is_same<half, T>::value>::type* = nullptr>
__device__ __forceinline__ void FastSpecializedAtomicAdd(T* base, size_t offset,
const size_t length, T value) {
atomicAdd(base + offset, value);
}
template<class T>
__device__ __forceinline__ void FastAdd(T* base, size_t offset, const size_t length, T value) {
FastSpecializedAtomicAdd(base, offset, length, value);两个 FastSpecializedAtomicAdd 重载靠 SFINAE 区分：std::is_same<half, T> 为真走上面那个（有三分支），否则透传给普通 atomicAdd——所以这个函数对 float、double 也能用，只是不做任何花活。FastAdd 只是给了一个更短的名字。
__global__ void dot(half* a, half* b, half* c, int n){
const int nStep = gridDim.x * blockDim.x;
half temp = 0.0;
int gid = blockIdx.x * blockDim.x + threadIdx.x;
while (gid < n) {
temp = temp + a[gid] * b[gid];
gid += nStep;
}
// atomicAdd(c, temp);
FastAdd(c, 0, 2, temp);
}kernel 和 v0 几乎一样，唯一区别是最后那行从 atomicAdd(c, temp) 换成了 FastAdd(c, 0, 2, temp)。第 66 行那句被注释掉的 atomicAdd(c, temp) 是原作者留的对照，方便切换两种实现。
“fastAtomicAdd”这个名字容易让人以为是手写的比较交换循环。在 sm_70 及以上、CUDA 10 以上的环境下（本机 sm_120 / CUDA 13.4 就在这个范围内），三个分支最终都直接调用硬件 atomicAdd，没有任何手写的 CAS。
代码顶部那行 #if CUDA_VERSION < 10000 || __CUDA_ARCH__ < 700 才会整段退化到 2 字节分支；而常说的 atomicCAS 循环模拟只存在于 sm_60 以下的 cuda_fp16.hpp 实现里，本机不经过。
这份实现是移植自 PyTorch 的 aten/src/ATen/native/cuda/KernelUtils.cuh。
完整代码
三个版本都很短，这里全部贴出。v0 和 v1 用 .cpp 版本，v2 用仓库的正式实现 fast_atomic_add_half.cu——原因是 v2/fast_atomic_add_half.cpp 那份工作副本本身是残缺的：它在 FastSpecializedAtomicAdd 的非 half 重载处截断，没有 FastAdd、没有 dot kernel、也没有 main，而且中间有个拼写错误 atmonicAdd，无法编译。
// ============================================================================
// atomic_add_half (v0) —— 基线：对每个 block 的部分和用纯 __half 的 atomicAdd
// ============================================================================
#include <cuda.h>
#include <cuda_runtime.h>
#include <time.h> // 【time.h 其实没用到：main 里没有计时代码，耗时靠 ncu 外部测】
#include <algorithm>
#include <iostream>
#include <cuda_fp16.h>
using namespace std;
#define N 32*1024*1024 // 数组长度：33,554,432 个 half
#define kBlockSize 256
// 【函数功能】两阶段内积 c = Σ a[i]*b[i]：
// 阶段① 每线程把分给自己的 a[gid]*b[gid] 累进寄存器 temp；
// 阶段② 每线程把自己那份 temp 用一条 atomicAdd 汇入同一个地址 c。
__global__ void dot(half* a, half* b, half* c, int n) {
const int nStep = gridDim.x * blockDim.x; //【grid-stride 步长】nStep = 总线程数 = N
half temp = 0.0; // 【线程私有的部分和】从 half 常量 0.0 起步
int gid = blockIdx.x * blockDim.x + threadIdx.x;
while (gid < n) {
temp = temp + a[gid] * b[gid];
gid += nStep; // 跳到下一个归本线程的元素
}
atomicAdd(c, temp);// 【本版的主角】2 字节宽的原子加：33,554,432 个线程全往 c 这一个地址发 atom.add.noftz.f16
}
int main(){
half *x_host = (half*)malloc(N * sizeof(half));
half *x_device;
cudaMalloc((void **)&x_device, N*sizeof(half));
for (int i = 0; i < N; i ++) x_host[i] = 0.1;
cudaMemcpy(x_device, x_host, N*sizeof(half), cudaMemcpyHostToDevice);
half *y_host = (half*)malloc(N * sizeof(half));
half *y_device;
cudaMalloc((void **)&y_device, N*sizeof(half));
for (int i = 0;i < N; i ++) y_host[i] = 0.1;
cudaMemcpy(y_device, y_host, N*sizeof(half), cudaMemcpyHostToDevice);
half *output_host = (half*)malloc(sizeof(half));
half *output_device;
cudaMalloc((void **)&output_device, sizeof(half));
cudaMemset(output_device, 0, sizeof(half));
// 【网格配置】block_num = (N+255)/256 = 131072，线程总数 = N，
// 与 v1/v2 完全一致（README 第 6 行：长度和 grid/block 全部对齐才好对比）。
int32_t block_num = (N + kBlockSize - 1) / kBlockSize;
dim3 grid(block_num, 1);
dim3 block(kBlockSize, 1);
dot<<<grid, block>>>(x_device, y_device, output_device, N); // 实参对形参：a←x_device、b←y_device、c←output_device（就那 1 个 half 槽）、n←N
cudaMemcpy(output_host, output_device, sizeof(half), cudaMemcpyDeviceToHost);
printf("%.6f\n", static_cast<double>(output_host[0])); // 【打印】实测 32.000000：增量 ~0.01，涨到 32 后 ulp=0.03125，加法被舍回原值（见文件头"坑 2"）
free(x_host);
free(y_host);
free(output_host);
cudaFree(x_device);
cudaFree(y_device);
cudaFree(output_device);
}
// ============================================================================
// atomic_add_half_pack2 (v1) —— 手写 pack：把两个相邻 half 打包成 __half2，
// 最后一次 4 字节宽的 atomicAdd
// ============================================================================
//相比v0改的只有两点：
// 1) 访存视图：不再把数组当成 half*，而是当成 Pack<half,2>*（两个 half 一组），
// 线程 i 处理第 i 组的 2 个元素，所以循环上界变成 n / pack_size；
// 2) 原子加：不再 atomicAdd(half*)，而是把 temp 塞进 half2 的两个 lane
// （x 和 y 都放同一个值），一次 atomicAdd(__half2*)，一次改 4 字节。
// v0 每次原子操作只搬 2 字节，占不满按 32-bit 粒度组织的访存/原子通路；
// 这里次数不变、每次宽度翻倍，A100 实测从 422.36ms 降到 137.02ms（README 表）。
#include <cuda.h>
#include <cuda_runtime.h>
#include <time.h>
#include <algorithm>
#include <iostream>
#include <cuda_fp16.h>
using namespace std;
#define N 32*1024*1024 // 元素个数（half 计），必须是 pack_size 的偶数倍，见文件头
#define kBlockSize 256
// 【pack 类型】alignas(4) 是灵魂：没有它，Pack<half,2> 只按成员最大对齐(2)，
// 数组排布就可能落在奇数地址上，后面 reinterpret_cast 成 half2* 就破坏对齐。
template<typename T, size_t pack_size>
struct alignas(sizeof(T) * pack_size) Pack {
T elem[pack_size];
};
template<typename T, int32_t pack_size>
__device__ __inline__ void AtomicAdd(Pack<T, pack_size>* address,
T val) {
#pragma unroll
for (int i = 0; i < pack_size; ++i) {
atomicAdd(reinterpret_cast<T*>(address) + i, static_cast<T>(val));
}
}
template<>
__device__ __inline__ void AtomicAdd<half, 2>(Pack<half, 2>* address, half val) {
half2 h2_val;
h2_val.x = static_cast<half>(val);
h2_val.y = static_cast<half>(val);
atomicAdd(reinterpret_cast<half2*>(address), h2_val);
}
template<typename T, int32_t pack_size>
__global__ void dot(Pack<T, pack_size>* a, Pack<T, pack_size>* b, Pack<T, pack_size>* c, int n){ // 实例化：T=half、pack_size=2（见 main 里的 dot<half, 2><<<...>>>）
const int nStep = gridDim.x * blockDim.x; // 线程总数 = N，组数 = N/2 → 多数线程空转一圈
T temp = 0.0;
int gid = blockIdx.x * blockDim.x + threadIdx.x;
while (gid < n / pack_size) { // 只遍历组；n 为奇数时会漏掉末尾元素（见文件头"长度对齐"）
for (int i = 0; i < pack_size; i++) { // 【一组 2 个元素都乘进来】a[gid].elem[i] 就是原来第 2*gid+i 个 half
temp = temp + a[gid].elem[i] * b[gid].elem[i];
}
gid += nStep;
}
AtomicAdd<T, pack_size>(c, temp); // 【每个线程一次 half2 原子加】后一半线程在这里加了个 0，次数与 v0 持平
}
int main(){ // 【输入 a 块】分配仍是 N 个 half，视图换成 Pack<half,2>*
half *x_host = (half*)malloc(N*sizeof(half));
half *x_device;
cudaMalloc((void **)&x_device, N*sizeof(half)); // cudaMalloc 的对齐足够放 half2
for (int i = 0; i < N; i++) x_host[i] = 0.1;
cudaMemcpy(x_device, x_host, N*sizeof(half), cudaMemcpyHostToDevice);
Pack<half, 2>* x_pack = reinterpret_cast<Pack<half, 2>*>(x_device); // 同一段内存，换类型看
// ---------- 输入 b：同上 ----------
half *y_host = (half*)malloc(N*sizeof(half));
half *y_device;
cudaMalloc((void **)&y_device, N*sizeof(half));
for (int i = 0; i < N; i++) y_host[i] = 0.1;
cudaMemcpy(y_device, y_host, N*sizeof(half), cudaMemcpyHostToDevice);
Pack<half, 2>* y_pack = reinterpret_cast<Pack<half, 2>*>(y_device);
// ---------- 输出 c：必须开 2 个 half ----------
// 一次 half2 原子加写 4 字节；只分配 2 字节会越界（v0 与这里的区别）。
half *output_host = (half*)malloc(2 * sizeof(half));
half *output_device;
cudaMalloc((void **)&output_device, 2 * sizeof(half));
cudaMemset(output_device, 0, sizeof(half) * 2); // 两个 lane 都清 0
Pack<half, 2>* output_pack = reinterpret_cast<Pack<half, 2>*>(output_device);
// 【网格配置与 v0 逐字相同】线程数= N 保持不变，才谈得上"只改原子加方式"的对比。
int32_t block_num = (N + kBlockSize - 1) / kBlockSize;
dim3 grid(block_num, 1);
dim3 block(kBlockSize, 1);
dot<half, 2><<<grid, block>>>(x_pack, y_pack, output_pack, N); // 实参对形参：a←x_pack、b←y_pack、c←output_pack（1 组=4 字节）、n←N（half 总数）
cudaMemcpy(output_host, output_device, 2 * sizeof(half), cudaMemcpyDeviceToHost);
printf("%.6f\n", static_cast<double>(output_host[0])); // 【打印 x lane】实测 64.000000：增量 ~0.02，涨到 64 后 ulp=0.0625，加法被舍回原值（见 ../v0/ 文件头"坑 2"）
free(x_host);
free(y_host);
free(output_host);
cudaFree(x_device);
cudaFree(y_device);
cudaFree(output_device);
}#include <cuda.h>
#include <cuda_runtime.h>
#include <time.h>
#include <algorithm>
#include <iostream>
#include <cstdint>
#include <cuda_fp16.h>
using namespace std;
#define N 32*1024*1024
#define kBlockSize 256
// FastAdd is referenced from
// https://github.com/pytorch/pytorch/blob/396c3b1d88d7624938a2bb0b287f2a19f1e89bb4/aten/src/ATen/native/cuda/KernelUtils.cuh#L29
template<typename T, typename std::enable_if<std::is_same<half, T>::value>::type* = nullptr>
__device__ __forceinline__ void FastSpecializedAtomicAdd(T* base, size_t offset,
const size_t length, T value) {
#if ((defined(CUDA_VERSION) && (CUDA_VERSION < 10000)) \
|| (defined(__CUDA_ARCH__) && (__CUDA_ARCH__ < 700)))
atomicAdd(reinterpret_cast<half*>(base) + offset, static_cast<half>(value));
#else
// Accounts for the chance base falls on an odd 16 bit alignment (ie, not 32 bit aligned)
__half* target_addr = reinterpret_cast<__half*>(base + offset);
bool low_byte = (reinterpret_cast<std::uintptr_t>(target_addr) % sizeof(__half2) == 0);
if (low_byte && offset < (length - 1)) {
__half2 value2;
value2.x = value;
value2.y = __float2half_rz(0);
atomicAdd(reinterpret_cast<__half2*>(target_addr), value2);
} else if (!low_byte && offset > 0) {
__half2 value2;
value2.x = __float2half_rz(0);
value2.y = value;
atomicAdd(reinterpret_cast<__half2*>(target_addr - 1), value2);
} else {
atomicAdd(reinterpret_cast<__half*>(base) + offset, static_cast<__half>(value));
}
#endif
}
template<typename T, typename std::enable_if<!std::is_same<half, T>::value>::type* = nullptr>
__device__ __forceinline__ void FastSpecializedAtomicAdd(T* base, size_t offset,
const size_t length, T value) {
atomicAdd(base + offset, value);
}
template<class T>
__device__ __forceinline__ void FastAdd(T* base, size_t offset, const size_t length, T value) {
FastSpecializedAtomicAdd(base, offset, length, value);
}
// vector inner product
__global__ void dot(half* a, half* b, half* c, int n){
const int nStep = gridDim.x * blockDim.x;
half temp = 0.0;
int gid = blockIdx.x * blockDim.x + threadIdx.x;
while (gid < n) {
temp = temp + a[gid] * b[gid];
gid += nStep;
}
// atomicAdd(c, temp);
FastAdd(c, 0, 2, temp);
}
int main(){
half *x_host = (half*)malloc(N*sizeof(half));
half *x_device;
cudaMalloc((void **)&x_device, N*sizeof(half));
for (int i = 0; i < N; i++) x_host[i] = 0.1;
cudaMemcpy(x_device, x_host, N*sizeof(half), cudaMemcpyHostToDevice);
half *y_host = (half*)malloc(N*sizeof(half));
half *y_device;
cudaMalloc((void **)&y_device, N*sizeof(half));
for (int i = 0; i < N; i++) y_host[i] = 0.1;
cudaMemcpy(y_device, y_host, N*sizeof(half), cudaMemcpyHostToDevice);
half *output_host = (half*)malloc(sizeof(half) * 2);
half *output_device;
cudaMalloc((void **)&output_device, sizeof(half) * 2);
cudaMemset(output_device, 0, sizeof(half) * 2);
int32_t block_num = (N + kBlockSize - 1) / kBlockSize;
dim3 grid(block_num, 1);
dim3 block(kBlockSize, 1);
dot<<<grid, block>>>(x_device, y_device, output_device, N);
cudaMemcpy(output_host, output_device, sizeof(half) * 2, cudaMemcpyDeviceToHost);
printf("%.6f\n", static_cast<double>(output_host[0]));
free(x_host);
free(y_host);
free(output_host);
cudaFree(x_device);
cudaFree(y_device);
cudaFree(output_device);
}
v0 的文件头里写得很直白：#include <time.h> 其实没用到，“main 里没有计时代码，耗时靠 ncu 外部测”。本文表格里的毫秒数是另加 cudaEvent 测的。
另外三份 main 都只 printf 出结果值，不与任何参考值比对——因为它们自己也知道结果是 32 而不是真值。
实测性能对比分析
规模 N = 33,554,432，线程数 = N，每个线程发一条原子指令，全部指向同一个地址。跑 7 次、取后 5 次最小值。这个负载的耗时完全由原子通路决定，带宽指标没有意义，所以下面换算成“每一次原子操作花了多少纳秒”来比较。
本机实测（RTX 5070）
| 版本 | 做法 | 原子宽度 | 耗时 (ms) | 每次原子 (ns) | 相对 v0 | 打印值 |
|---|---|---|---|---|---|---|
| v0 | half 原子加基线 | 2 字节 | 48.074 | 1.43 | 1.00× | 32.0 |
| v1 | 手写 pack 成 half2 | 4 字节 | 39.457 | 1.18 | 1.22× | 64.0 |
| v2 | fastAtomicAdd 运行期对齐 | 4 字节 | 39.456 | 1.18 | 1.22× | 32.0 |
三点观察。
第一，v1 和 v2 的耗时几乎一模一样（39.457 vs 39.456 ms）。这符合预期：两者最终的原子指令完全相同（4 字节 atom.add.noftz.f16x2），指令条数也相同，只是“怎么凑齐这 4 字节”的路径不同——v1 靠改数据布局，v2 靠运行期判断。两条路径的结果一致，所以性能一致。
第二，宽度翻倍只换来 1.22 倍，远不是 2 倍。如果原子通路的成本纯粹由字节数决定，2 字节变 4 字节应该接近 2 倍。实际只有 1.22 倍，说明这条通路上还有相当一部分成本是与宽度无关的固定开销——同地址操作排队、L2 往返延迟、指令发射。
第三，v2 的打印值 32 和 v1 的 64 不同，但这不是性能差异。v1 每次原子加贡献的是两个乘积之和（0.02），v2 因为走分支 1，只有本槽真正加到了 0.01、右邻居加的是 0——所以 v2 的累加速率与 v0 相同，冻结点也就和 v0 一样落在 32。这个差异纯粹是第 ① 段那个 half 精度问题的表现。
和 A100 上的数据对照
仓库 README 里原作者在 A100-PCIE-40GB 上用 ncu 测的同一组实验：
| 版本 | 原子宽度 | A100 耗时 (ms) | A100 每次原子 (ns) | A100 相对 v0 |
|---|---|---|---|---|
| v0 | 2 字节 | 422.36 | 12.59 | 1.00× |
| v1 | 4 字节 | 137.02 | 4.08 | 3.08× |
| v2 | 4 字节 | 137.01 | 4.08 | 3.08× |
把两张表对齐看，反差非常明显：
两个原因。一是 Blackwell 这一代对原子操作的硬件支持强得多——A100 是 Ampere，同地址原子操作的排队成本在它上面要重得多；本机 v0 的 1.43 ns/次，已经接近“L2 往返延迟被大量并发掩盖掉”之后的理论下限了。
二是“可优化的空间”本来就不一样。A100 上 2 字节原子操作的代价是 4 字节的 3 倍，说明那条通路在窄宽度下效率损失极大；本机只差 1.22 倍，说明 2 字节操作在这个架构上损失小得多。收益的上限由“基线有多差”决定，而不是由优化手法有多妙决定——这一点在 elementwise 那篇里也出现过同样的规律。
如果只看 A100 的表，很容易得出“把 half 原子加打包成 half2 能提速 3 倍”这个结论，然后在别的卡上照做、发现只有 1.22 倍而怀疑自己写错了。两张表放在一起才能看出：3.08× 和 1.22× 都不是代码的性质，是代码与硬件配对的结果。
小结
这三个版本把“凑宽度”这个思路的两条不同路径摆在一起，很值得对照：
v1 改数据结构：把数组重新解释成 Pack<half,2>*，让相邻两个 half 天然待在同一个 4 字节槽里。代价是对调用方提出要求——长度必须偶数、地址必须 4 字节对齐。做法简单，但不通用。
v2 该调用点：数据布局一点不动，只在实际发起原子操作时判断对齐，能凑成就借邻居的格子（给它加 0，不改变它的值），凑不成退回 2 字节。代价是每次多几条判断指令，换来的是任意长度、任意地址都能用——所以它能被抽成通用库函数（PyTorch 就是这么用的）。
两者的最终原子指令完全相同，实测耗时也几乎相同（39.457 vs 39.456 ms）。差别只在于“这份复杂度由谁来承担”。
另外一个更基础的收获是关于基线的：同一份优化，在 A100 上值 3.08 倍，在本机值 1.22 倍。原因不是代码变差了，而是本机的 2 字节原子加本来就不那么差。这和前面 elementwise（A100 上 1.69× / 本机 1.06×）、reduce（A100 上 9.94× / 本机 2.20×）是同一个规律：
看到一个 3 倍的优化案例，先问一句：它的基线为什么那么慢？
如果基线慢是因为硬件本身的短板（A100 的窄原子通路、A100 上 2 字节访存喂不饱 1555 GB/s），那换一张没有这个短板的卡，收益就会大幅缩水——甚至像 reduce 的 v1 那样变成负收益。
反过来，真正可移植的是“先定位瓶颈，再决定要不要动手”这个方法本身。
这一篇也是五篇里的最后一篇。五个算子走下来，带宽受限（reduce、elementwise、upsample、gemv）和原子竞争（本篇）这两类问题各占一半，而每一类的优化空间都高度依赖目标硬件——这大概是这组实测里最一致的一条结论。