02 · CUDA from a Systems Engineer’s Eyes · 报告 03
仍是同一块 RTX 4090。归约没有复用可做：每个元素只读一次， 所以上限一秒钟就能算出来：4N ÷ 1008 GB/s。 于是这一章只有一张成绩单：一个百分比。含糊的地方都会暴露。 书按顺序给了五个 kernel，每一步都有一套漂亮的分析。 开跑前，我先写死了五条预测和判决标准。 结果：五步优化只有一步打中了瓶颈， 它让每个元素快 292 倍；其余四步的收益是零。
N 元素的求和归约只有一件必须做的事：把 N 个 float 从显存读进来。 读 4N 字节，只读一次，不可能更少。写只有原子加： 次数等于 grid 数，每次 4 字节，和 1 GiB 相比可以忽略。
下限时间 = 4N 字节 ÷ 1008 GB/s N = 2²⁸（1 GiB） → 1,073,741,824 ÷ 1008e9 = 1.0652 ms N = 任意 → 永远是这条线 // 报告 01/02 那套"算术强度 → roofline"的推导，在这里退化成一条水平线。 // 省掉了所有争论 —— 但也意味着任何"减少指令条数"的优化都动不了它。
所以这一章唯一有意义的成绩单是一个百分比：下限时间 ÷ 实测时间（达成率）。 这个下限简单到在纸上一秒就能算出来， 所以含糊的地方都会暴露。
浮点运算数 = N − 1 // 归约树总操作数，和串行一样 必须搬的字节 = 4N // 每个 float 读一次，就这一次 算术强度 = (N−1) / 4N ≈ 0.25 FLOP/B 脊点 = 82.6 TFLOP/s ÷ 1008 GB/s = 83.4 FLOP/B 0.25 ÷ 83.4 = 1/334 // 比脊点低两个半数量级 // 对比这一路走来的几个 kernel（报告 02 §1 主表） K1 朴素 SGEMM 0.250 FLOP/B // 巧了，一样 K3 分块 SGEMM 7.969 K5 二维粗化 SGEMM 12.720 归约（任何版本） 0.250 —— 而且没有任何办法提高它
| 这一步 | 做了什么 | 书里衡量的 | 实际收益 | 说明 |
|---|---|---|---|---|
| R1 → R2 | 消除控制 / 访存发散 | 资源利用率 29.5% → 66.4% | 微秒级 | 定性有效，定量不可信 |
| R2 → R3 | 搬进 shared memory | 显存访问 ~N·log N → N+1 | 零 | 8 KB 工作集全在 L1 里；收益甚至略负 |
| R3 → R4 | 分段 + 原子加 | 书里只说"能处理任意长度" | 每元素快 292 倍 | 唯一打中瓶颈的一步 |
| R4 → R5 | 线程粗化 | __syncthreads() 次数与原子加都除以 4 | 零 | 1.1203 → 1.1211 ms |
| R5 → R6 | warp shuffle（书里没有） | 最后 5 轮不走 shared memory | 零 | 1.1211 → 1.1212 ms |
判决标准在开跑前定死，所以 §8·二 那一栏没有商量余地。 这个做法直接沿用报告 02 §3 的"先算上限、再测、再对账"。
| # | 预测 | 判决标准 | 理由 |
|---|---|---|---|
| ① | R1 → R2 有明显加速 | 净耗时之比 > 2× | 两者都只用 1 个 SM，但请求条数确实少了 |
| ② | R2 → R3 再快一截 | 净耗时之比 > 1.5× | 显存访问降到 N+1，向下限靠近 |
| ③ | R3 → R4 是数量级的跳变 | 每元素耗时之比 > 50× | 本章的主赌注 —— 第一次动到真正的瓶颈 |
| ④ | R4 → R5 收益很小 | 加速 < 1.2× | 字节账没变，收益只能来自启动开销和原子加 |
| ⑤ | R5 → R6 收益更小 | 加速 < 1.1× | 如果 R5 已经到八成带宽，这一步基本看不见 |
赌注一句话
真正动瓶颈的只有 R3 → R4 那一步，其余四步的收益会在 4090 上迅速收敛到零。 书里的分析每一条都对。它衡量的是 一个 SM 内部的执行资源，而这在这张卡上不是最紧的约束。
// ---------------------------------------------------------------------
__global__ void reduce_r1(float *input, float *output) {
unsigned int i = 2 * threadIdx.x;
for (unsigned int stride = 1; stride <= blockDim.x; stride *= 2) {
if (threadIdx.x % stride == 0) // 编号是 stride 倍数的才干活
input[i] += input[i + stride]; // → warp 内部分裂 = 控制发散
__syncthreads(); // 必须在 if 外面
}
if (threadIdx.x == 0) *output = input[0];
}
R1 的设计遵循"owner computes"：每个位置只归一个线程所有，只有它能写。
线程 t 认领 input[2t]，
即偶数位置。stride 从 1 开始翻倍，
条件 threadIdx.x % stride == 0 挑出这一轮的活跃线程。
for 循环的一轮；深色格子是这一轮被写入的位置。
第 1 轮偶数位置拿到成对的和（0–1、2–3…），
第 2 轮下标是 4 的倍数的位置拿到四个相邻元素的和（0–3、4–7…）。
问题在于：owner 之间越来越远，
活跃线程撒在每个 warp 里。
"执行资源利用率"衡量的是：硬件腾出的线程槽位里，有多少真的产出了结果。 它是一个分数：分子是产出，分母是占用。 下面分别数这两个量。
分母的关键：占用以 warp 为单位，不以线程为单位
一个 warp 只要还有一个线程活着，就占满 32 个槽位。 原因在 SIMD 硬件：32 个 lane 步调一致， 不干活的 lane 只是被掩码挡住，槽位照样占着。 所以分母数的是"活跃 warp 数"， 不是"活跃线程数"。这一条是整个计算的支点。
取 N = 256：启动 N/2 = 128 个线程，即 4 个 warp（warp w 装 tid 32w … 32w+31）。
stride 从 1 翻倍到 128，共 8 轮；
每轮由 tid % stride == 0 挑出活跃线程。
stride 1 2 4 8 16 32 64 128 活跃线程 128 64 32 16 8 4 2 1 ← tid % stride == 0 每 warp 几个 32 16 8 4 2 1 · · 活跃 warp 4 4 4 4 4 4 2 1 ← 有一个活的就算一个 └────────── 前 6 轮 4 个 warp 全在场 ──────────┘ // 要 stride 走到 64，才第一次出现"整个 warp 一个活跃线程都没有"。
分子 · 真正提交的结果 128 + 64 + 32 + 16 + 8 + 4 + 2 + 1 = 255 // = N − 1。把 256 个数归约成 1 个，本来就需要 255 次加法 —— 串行也是这个数。 分母 · 被占住的线程槽位 (4 + 4 + 4 + 4 + 4 + 4 + 2 + 1) × 32 = 27 × 32 = 864 // 27 是"活跃 warp 数"在 8 轮里的总和，×32 是每个 warp 占满一整组槽位。 利用率 = 255 ÷ 864 = 29.5% // 硬件掏出去的槽位里，只有不到三成产出了一次加法。
浪费从哪来：分子固定，分母被"陪跑"撑大
分子 255 是任何归约实现都躲不掉的下限，它不会变。
所以要提高利用率，只能压低分母。
分母是 864，因为前 6 轮里每一个 warp 都至少留着一个活跃线程：
第 6 轮只有 4 个线程在干活，却仍占着 4×32 = 128 个槽位。
根源在 owner 的摆法：
2t 让活跃线程均匀撒在每个 warp 里，哪个 warp 都退不了休。
下一节的 R2 只改了三处：分子一个数都没动，分母从 864 降到 384， 利用率从 29.5% 升到 66.4%。
// ---------------------------------------------------------------------
__global__ void reduce_r2(float *input, float *output) {
unsigned int i = threadIdx.x;
for (unsigned int stride = blockDim.x; stride >= 1; stride /= 2) {
if (threadIdx.x < stride) // 活跃的永远是编号最小的一批
input[i] += input[i + stride];
__syncthreads();
}
if (threadIdx.x == 0) *output = input[0];
}
owner 从 2*threadIdx.x 换成 threadIdx.x，
stride 从 blockDim.x 开始减半，条件从 % 换成 <。
加法次数不变，
变的只是"谁来干"。
仍取 N = 256：128 个线程，即 4 个 warp；stride 从 128 减半到 1，共 8 轮。
下表照 §2·三 的表再列一遍，
末尾多加一行。两本账都能从这张表里读出来。
stride 128 64 32 16 8 4 2 1 活跃线程 128 64 32 16 8 4 2 1 ← tid < stride 活跃 warp 4 2 1 1 1 1 1 1 ← 整批退休 每条访问几条请求 1 1 1 1 1 1 1 1 ← 地址连号，永远只落在一条线上 // 对比 R1 的同两行：活跃 warp 4 4 4 4 4 4 2 1 // 每条访问 2 2 2 2 2 1 1 1 ← 为什么是 2，见下面那一栏
"每 warp 几条请求"这一行是什么意思
一条 load / store 是整个 warp 一起发出的一条指令。硬件拿到 每个活跃 lane 各自算出的地址，最多 32 个； 这些地址落在几条 128 字节的线上，就发几条请求。 判据只有一个：这组地址覆盖了几条线。 地址连号，就挤在一条线里；地址分散，就要多取几条，每条里还有用不上的字节。
// 一"线" = 128 字节 = 32 个 float。下标和线号的换算： 线 0 = input[0…31] 线 1 = input[32…63] 线 2 = input[64…95] 线 4 = input[128…159] … // 注意 input[i] += input[i + stride] 是【读-改-写】，一轮里其实有三条访问， // 每条各有自己的一组地址，各判各的。 R1 stride = 1，warp 0（tid 0…31，i = 2·tid） 读 input[i] 地址 0, 2, … 62 分散 → 线 0 + 线 1 2 条 读 input[i+1] 地址 1, 3, … 63 分散 → 线 0 + 线 1 2 条 写 input[i] 地址 0, 2, … 62 分散 → 线 0 + 线 1 2 条 // 三条都要取回 256 字节，而真正用得上的只有 128 字节 —— 一半白搬 R2 stride = 128，warp 0（tid 0…31，i = tid） 读 input[i] 地址 0, 1, … 31 连号 → 线 0 1 条 读 input[i+128] 地址 128, … 159 连号 → 线 4 1 条 写 input[i] 地址 0, 1, … 31 连号 → 线 0 1 条 // 三条各取回 128 字节，一个字节都不浪费。地址落在哪条线不要紧，要紧的是"只落在一条"
所以这一行只由"活跃 lane 这一刻请求的地址"决定， 和它们平时"拥有"哪些位置无关。 活跃 lane 变少时这个数变不变， 要看剩下的地址还跨不跨线。R1 一直跨线，直到只剩 1 个 lane 才降到 1 条； 下面逐轮看。
账一 · 执行资源利用率 分子 = 128 + 64 + … + 1 = 255 // 和 R1 一模一样，加法一次没少 分母 = (4 + 2 + 1 + 1 + 1 + 1 + 1 + 1) × 32 = 12 × 32 = 384 利用率 = 255 ÷ 384 = 66.4% // R1 是 255 ÷ 864 = 29.5% 账二 · 发给显存的请求条数 每个活跃线程每轮 2 读 + 1 写 = 3 次访问，每次访问按上表折算成请求： R2 = (4 + 2 + 1 + 1 + 1 + 1 + 1 + 1) × 1 × 3 = 12 × 3 = 36 R1 = (4×2 + 4×2 + 4×2 + 4×2 + 4×2 + 4 + 2 + 1) × 3 = 47 × 3 = 141 141 ÷ 36 = 3.9 倍
逐轮看一遍 R1：为什么中间几轮还是 2，到 stride = 32 才变 1
每一轮，warp 0 里活着的 lane 各请求一个地址；下表看这组地址落在哪几条线上。
活跃 lane 由 tid % stride == 0 挑出。
下表以三条访问里的第一条（读 input[i]）为例，
另外两条结论相同。
stride warp 0 还活着的 tid 它们 owner 的下标 落在哪几条线 请求 1 0,1,…,31（32 个） 0,2,…,62 线0 + 线1 2 2 0,2,…,30（16 个） 0,4,…,60 线0 + 线1 2 4 0,4,…,28（8 个） 0,8,…,56 线0 + 线1 2 8 0,8,16,24（4 个） 0,16,32,48 线0 + 线1 2 16 0,16（2 个） 0,32 线0 + 线1 2 32 0（1 个） 0 只剩线0 1 // 关键：活跃 lane 越来越少，但剩下的那几个地址仍然【一头一尾地跨着两条线】—— // 最小的永远是 input[0]（线0）；最大的只要还有第二个 lane 活着，就一定 ≥ input[32]（线1）。 // 所以 lane 数从 32 掉到 2，请求数一直是 2。直到只剩 tid 0 一个人，才塌回 1 条。 // 这就是 2 2 2 2 2 1 1 1 的由来 —— 变的是"跨不跨线"，不是"有几个 lane"。 // 一轮里其实有三条访问（读 input[i]、读 input[i+stride]、写 input[i]），各判各的， // 但三条的结论在这里完全一致，所以上面那笔账才能直接乘 3。
R2 那一行全是 1，原因一句话就能说清：
活跃 lane 是 tid < stride
挑出的一段连号，请求的地址就是 input[0…stride−1]，
最宽 32 个 float = 128 字节，正好一条线。
lane 变少只会让这段更短，永远跨不到第二条线。
⚠ 要当心的是模型本身，不是这套算术。 现代 GPU 的 L1/L2 按 32 字节 sector 计数， 和"一块 128 字节 = 一条请求"不是一一对应。方向是对的，但倍数不要当成硬件上的精确值。
N 越大，两边差得越多。 R1 的浪费正比于 warp 总数；R2 的浪费是常数， 因为只有尾巴上那几轮剩一个 warp 在发散。
N = 2048（1024 线程 = 32 个 warp） 利用率 R1 2047 ÷ 7136 = 28.7% R2 2047 ÷ 2176 = 94.1% 请求数 R1 383 × 3 = 1149 R2 68 × 3 = 204 1149 ÷ 204 = 5.6 倍 // R2 剩下的那点浪费，全在"活跃线程掉到 32 以下之后"的最后几轮 —— // 整棵树困在一个 warp 里，R2 治不了。那正是 R6 要收拾的尾巴。
// ---------------------------------------------------------------------
__global__ void reduce_r3(const float *input, float *output) {
extern __shared__ float s[];
unsigned int t = threadIdx.x;
s[t] = input[t] + input[t + blockDim.x]; // 第 1 轮 + 搬进 shared
for (unsigned int stride = blockDim.x / 2; stride >= 1; stride /= 2) {
__syncthreads(); // 挪到循环头：同步上一行那次写
if (t < stride) s[t] += s[t + stride];
}
if (t == 0) *output = s[0];
}
blockDim.x/2 起步，__syncthreads() 也挪到了循环头。
另一个好处：R3 不再改写 input。
R1/R2 都是原地归约，跑完输入就毁了。
R3 访问 global memory 的地方只剩三处，请求账很好算。 三处是开头的两条读和结尾的一次写，中间的归约树全在片上。 数法和 §3·三 相同。
// N = 256 → 128 个线程 = 4 个 warp。全局访问就这三条： 读 input[t] warp 0 的地址 0 … 31 连号 → 线 0 1 条 4 个 warp 各 1 条 4 条 读 input[t + blockDim.x] warp 0 的地址 128 … 159 连号 → 线 4 1 条 4 个 warp 各 1 条 4 条 写 output 只有 tid 0 发 1 条 ───────── 合计 9 条 // 换个角度数，同一个答案：N 个元素每个只读一次，连号 —— // N 个 float ÷ 32 个 float/线 = N/32 = 8 条线，再加 1 次写 = 9。 请求数 R2 = 36 → R3 = 9 又少 4 倍
请求数只剩 R2 的 1/4，省在哪里
不是每条请求变便宜了，而是数据不再反复进出 global memory。
R2 的 36 条请求里，每一轮归约都把部分和写回 global memory，
下一轮再读回来：8 轮、每轮 3 次访问，全部走 global memory。
R3 把第 1 轮和搬运合成一步，之后 s[] 一直在片上。
global memory 访问从"每轮都有"降到"只在开头和结尾"，即 N + 1 次。
字节账到此触底，后面三步一个字节也省不掉。
剧透：在 N=2048 上，这一步没有变快，反而略慢
8 KB 的工作集本来就全在 L1 里， R3 省掉的那些 global memory 访问一次也没到过显存。 书说 global memory 访问从 ~N·log N 降到 N+1，这没错。 但它数的是访问次数，不是显存流量。 完整分析见 §8·三。
// ---------------------------------------------------------------------
__global__ void reduce_r4(const float *input, float *output) {
extern __shared__ float s[];
unsigned int segment = 2 * blockDim.x * blockIdx.x; // 我这个 block 的段起点
unsigned int i = segment + threadIdx.x;
unsigned int t = threadIdx.x;
s[t] = input[i] + input[i + blockDim.x];
for (unsigned int stride = blockDim.x / 2; stride >= 1; stride /= 2) {
__syncthreads();
if (t < stride) s[t] += s[t + stride];
}
if (t == 0) atomicAdd(output, s[0]); // 每个 block 交一次答卷
}
2*blockDim.x 个元素。每个 block 独立跑完自己那一段的树，
再用 atomicAdd 把部分和加到同一个 output 上。
__syncthreads() 管不到 block 之间，所以只能分这两层。
__syncthreads() 只在 block 内有效，而一个 block 只能待在一个 SM 上。
R3：处理 2048 个元素，耗时 0.0025 ms → 每元素 1.2207 ns R4：处理 268435456 个元素，耗时 1.1203 ms → 每元素 0.0042 ns 数据量 ×131072 ←─ 十三万倍 耗时 ×448 ←─ 只多了四百多倍 每元素 快 292 倍 （用净耗时口径算也有 82 倍） // 128 个 SM 闲着的时候，你在那一个 SM 里做什么优化都是徒劳。
R4 一步把达成率推到 95.1%。达成率只有一个公式，有两种等价写法：
达成率 = 下限时间 ÷ 实测时间 = (4N ÷ 峰值带宽) ÷ 耗时 = 有效带宽 ÷ 峰值带宽 = (4N ÷ 耗时) ÷ 峰值带宽 // 两种写法里 4N 都在分子分母间约掉了 —— 本质上就是"你比理论下限慢了多少"。 代进 R4 的数（N = 2²⁸，耗时 1.1203 ms）： ① 先算有效带宽 4N = 4 × 268,435,456 = 1,073,741,824 字节 = 1 GiB 有效带宽 = 1,073,741,824 ÷ 0.001120 s = 958.5 GB/s 达成率 = 958.5 ÷ 1008 = 95.1% ② 先算下限时间（同一个数） 下限 = 1,073,741,824 ÷ 1008e9 = 1.0652 ms 达成率 = 1.0652 ÷ 1.1203 = 95.1%
这个"达成率"和报告 01/02 里那个不是一回事
这里的分母就是带宽本身。
报告 02 的 roofline 达成率，分母是
算术强度 × 带宽，是推出来的算力上限，所以会出现 2112% 这种被证伪的数。
归约的算术强度恒为 0.25，一步也动不了（§1·二），
所以 roofline 退化成一条水平线，分母简化成峰值带宽。
因此这一章的成绩单很干净：它没有可被证伪的"假设"。
R4 跑完，墙上已经没有空间留给 R5 和 R6 了。 1008 GB/s 是理论峰值。STREAM 类基准通常只能到 85–93%， 归约是纯顺序读，所以能更贴墙。
// ---------------------------------------------------------------------
__global__ void reduce_r5(const float *input, float *output) {
extern __shared__ float s[];
unsigned int segment = COARSE_FACTOR * 2 * blockDim.x * blockIdx.x;
unsigned int i = segment + threadIdx.x;
unsigned int t = threadIdx.x;
float sum = input[i]; // 粗化循环：全员活跃
for (unsigned int tile = 1; tile < COARSE_FACTOR * 2; ++tile)
sum += input[i + tile * blockDim.x]; // 步长 = blockDim.x → 合并
s[t] = sum;
for (unsigned int stride = blockDim.x / 2; stride >= 1; stride /= 2) {
__syncthreads();
if (t < stride) s[t] += s[t + stride];
}
if (t == 0) atomicAdd(output, s[0]);
}
sum 里，四个元素全加完才写进 shared memory，线程之间没有依赖。
⚠ 我们的 reduce_ch10.cu
用的是 4，是图里的 2 倍（每线程 8 个元素、粗化循环 7 步）。
__syncthreads() 和 shared memory 访问都减半。
书这套推理的前提是"硬件资源不够、block 会被串行"。
这个前提在 4090 上确实成立：R4 要跑 52 万个 block，同时只驻留得下几百个。
但实测收益还是零，
因为省下的是"算"的时间，而这个 kernel 的墙在"搬"那一侧。
字节账一分没变，这一步改的全是开销账。 实测是 R4 的 0.999×， 即收益为零（1.1203 → 1.1211 ms）。
按行号逐行读。注释里标了四个阶段：
① 粗化循环（和 R5 一字不差）→
② warp 内 5 轮 shuffle → ③ 每个 warp 交一个数 →
④ 第一个 warp 收尾。
新增的只有 ②③④，
它们只省下一件事：__syncthreads() 从 8 次降到 1 次。
01 __global__ void reduce_r6(const float *input, float *output) {
02 extern __shared__ float warp_s[]; // 只要 blockDim.x/32 个 float —— R5 要 blockDim.x 个
03 unsigned int segment = COARSE_FACTOR * 2 * blockDim.x * blockIdx.x;
04 unsigned int i = segment + threadIdx.x; // 03–05 行和 R5 一字不差
05 unsigned int t = threadIdx.x;
06 unsigned int lane = t & 31, warp = t >> 5; // 我是第几个 warp 的第几个 lane
07 unsigned int nWarps = blockDim.x >> 5; // BLOCK=256 → 8 个 warp
08
09 float sum = input[i]; // ① 粗化循环：和 R5 完全相同
10 for (unsigned int tile = 1; tile < COARSE_FACTOR * 2; ++tile)
11 sum += input[i + tile * blockDim.x]; // 这一段才是真正搬字节的地方
12
13 for (int off = 16; off > 0; off >>= 1) // ② warp 内 5 轮，全在寄存器
14 sum += __shfl_down_sync(0xffffffffu, sum, off);
15 // 做完：每个 warp 的 lane 0 手里是本 warp 32 个数的和
16 if (lane == 0) warp_s[warp] = sum; // ③ 8 个 warp 各交 1 个数
17 __syncthreads(); // 全 kernel 唯一一次屏障（R5 有 8 次）
18
19 if (warp == 0) { // ④ 头一个 warp 收尾 —— 32 个 lane 整齐进出，mask 才合法
20 sum = (lane < nWarps) ? warp_s[lane] : 0.0f; // 多出来的 lane 取单位元 0
21 for (int off = 16; off > 0; off >>= 1)
22 sum += __shfl_down_sync(0xffffffffu, sum, off);
23 if (lane == 0) atomicAdd(output, sum);
24 }
25 }
4 2 1 1 1 1 1 1，尾巴上一共六个 1：
从 stride = 32 起，整棵树只剩一个 warp。
第一个 1 那轮还是满员的 32 个 lane，不发散；
问题出在后五个：warp 内只剩 16、8、4、2、1 个 lane 在干活，
却每轮还要访问 shared memory、调一次 __syncthreads()。
R6 让这 5 轮改走寄存器；每个 warp 只交一个数给 shared memory，
再由 warp 0 收尾。全 kernel 只剩一次 __syncthreads()。
R6 是整条链上唯一一个书里没有的版本，也是唯一一个"理论上该有收益"的收尾优化。 实测收益同样为零：1.000× （1.1211 → 1.1212 ms）。 原因是 R4 已经站在带宽墙上， 省掉再多指令也没有地方可以变快。
收益是零，但这套原语值得单独学
__shfl_*_sync 是 vLLM 的
layernorm_kernels.cu、FlashAttention 的 warp 级归约里的标配。
它在这个 kernel 上没有收益，是因为这里的墙是带宽；
换一个不受带宽限制的算子，同一套写法就有实际收益。
下面按行号读一遍：
先分清几个数字，再看一条 shuffle 指令做了什么，
然后看收尾的 warp，最后和 R5 的 shared memory 树逐项对照。
第 06–07 行的三个位运算都基于一件事：32 = 2⁵。 "除以 32"就是右移 5 位，"模 32"就是保留低 5 位。 换句话说，线程编号的二进制本来就分成两段： 高位段是 warp 号，低 5 位是 lane 号。 这两行不是技巧， 只是读出这个数本来的结构。
// 拿 t = 77 摊开（BLOCK = 256） t = 0b …0100 1101 (77) 31 = 0b …0001 1111 ← 低 5 位全 1，高位全 0 t & 31 = 0b …0000 1101 (13) ← lane：我在本 warp 里排第 13 t >> 5 = 0b …0000 0010 (2) ← warp：我在第 2 个 warp（低 5 位整个移出去丢掉） t = 77 : 0000010 01101 验算：2 × 32 + 13 = 77 ✓ └──┬──┘ └─┬─┘ warp = 2 lane = 13
为什么这几个变量一律写成 unsigned int
对 unsigned，% 32 和 & 31
编译出的机器码完全一样——
写位运算只是表意习惯（"我在取 lane"比"我在取余数"直白）。
换成 signed int 就不等价了：
-1 % 32 == -1 而 -1 & 31 == 31
-1 / 32 == 0 而 -1 >> 5 == -1 // 算术右移向负无穷取整，除法向零取整
所以对 signed，编译器不能做这个替换，
必须插额外指令处理符号位。
这条对整份 CUDA 代码都成立：
从 threadIdx / blockIdx 派生的下标一律用 unsigned
——这样语义正确，这些除法和取模也真正免费。
这么切为什么是"对的"，以及多维 block 的经典 bug
硬件规则：把 block 内的线程按线性编号排好，
每 32 个连续编号切一刀。这个划分是固定的，改不了。
一维 block 的线性编号就是 threadIdx.x，
所以可以直接写 t & 31 / t >> 5。
多维 block 必须先线性化：
threadIdx.x + blockDim.x*threadIdx.y + blockDim.x*blockDim.y*threadIdx.z。
在二维 block 里直接把 threadIdx.x & 31 当 lane 是错的。
第 02 行的 extern 是 CUDA 借用的关键字，
和 C 里"符号定义在别的编译单元"的含义无关。
// 静态：大小写死在代码里，编译期确定 —— 书里 Fig. 10.13 / 10.15 的写法 __shared__ float input_s[BLOCK_DIM]; // BLOCK_DIM 必须是编译期常量 // 动态：大小留空，运行期才确定 —— 本报告 R3–R6 的写法 extern __shared__ float warp_s[]; // 注意数组没有维度：是 float[] 不是 float[8]
没有维度的数组是不完整类型。标准 C 只允许
extern 声明这样写，CUDA 借这个语法位置表达另一件事：
"这块 shared memory 有多大，不由 kernel 决定，由启动方在运行时给出。"
给出大小的地方是 <<< >>> 的第三个参数：
reduce_r6<<<grid, block, smem>>>(in, out); ^^^^ 动态 shared 的字节数，省略则为 0 // reduce_ch10.cu 的 launch 里给 R6 算的就是它： size_t smem = (kind == K_R6) ? (block.x / 32) * sizeof(float) // 8 个 float = 32 B : block.x * sizeof(float); // R5 要 256 个 = 1 KB // 单位是字节，不是元素个数 —— 两个高频错误： <<<grid, block, block.x / 32>>> // ✗ 漏了 sizeof，只给 8 字节（要 32） <<<grid, block>>> // ✗ 干脆忘了第三个参数 → 0 字节 → 第 16 行立即越界 // 两种都能编译、都能跑，只是结果错或随机崩。compute-sanitizer --tool memcheck 能抓到。
一个限制：整个 kernel 只有一块动态 shared memory
extern __shared__ float a[]; extern __shared__ int b[]; // a 和 b 是同一块内存，不是两块 // 要放多个数组，得自己切： extern __shared__ char smem[]; float* a = (float*)smem; int* b = (int*)(smem + n * sizeof(float));
数组名和类型只是一个"视角"
——数组名代表的总是那块内存的起始地址，类型只决定怎么解释里面的字节。
R6 用动态写法，是为了在命令行上换 BLOCK 而不必重编译；
代价是多了"忘传第三个参数"这个出错点。
难点多半在第 13–14 行。先记住一条：sum 不是一个变量，是 32 个变量。
第 09 行 float sum = input[i];
在每个线程里各执行一遍。一个 warp 有 32 个线程，就有 32 个互不相干的 sum，
各自在自己的寄存器里。它们同名，但不是同一个东西。
不先明确这一点，第 14 行出现两次的
sum 就读不通。
// off = 16 时，第 14 行展开后是这 32 条语句，【同时】发生（下标 = lane 号）： lane 0 : sum₀ = sum₀ + sum₁₆ ← 拿 lane 0+16 的值 lane 1 : sum₁ = sum₁ + sum₁₇ ⋮ lane 15 : sum₁₅ = sum₁₅ + sum₃₁ ← 以上 16 个有意义 ───────────────────────────────────── lane 16 : sum₁₆ = sum₁₆ + sum₁₆ ← 想拿 lane 32，越界 → 拿到自己 → 翻倍 ⋮ lane 31 : sum₃₁ = sum₃₁ + sum₃₁ ← 以上 16 个算出垃圾，但没人要 // __shfl_down_sync(0xffffffffu, sum, 16) 对 lane i 而言，返回值就是 sum(i+16)。 // 仅此而已 —— 它不改任何人的值，只是让 lane i 能读到 lane i+16 的寄存器。
三个参数各自的角色：var 最容易搞反
__shfl_down_sync( mask , var , delta )
│ │ └─ 我要往后数几个 lane 去取
│ └──────── 我摆出来给别人看的值（不是我要取的值）
└─────────────── 哪些 lane 在这里汇合，第 i 位对应 lane i
// 32 个 lane 各自摆出自己的 sum，然后各自按 delta 去取别人摆出来的那份。
越界的 lane 拿到的是它自己的值，不是 0，也不是未定义值。
i + delta ≥ 32 时源 lane 不会绕回，
CUDA 文档的说法是 "the upper delta lanes will remain unchanged"。
这个行为有明确定义，所以可以确认不会有 NaN 或陷阱值污染 lane 0 的结果。
⚠ mask 是一个承诺，不是筛选条件。
它的意思是"这些 lane 保证都会执行到这条指令"，
不是"只让这些 lane 干活"。承诺了 32 个、实际只到了 8 个，是未定义行为：不报错，换个架构结果就变。
所以第 19 行写
if (warp == 0) 是安全的（一个 warp 的 32 个 lane 要么全进、要么全不进）；
写成 if (t < 8) 就不安全，因为它把一个 warp 切碎了。
"lane 0 读 sum₁₆ 时，lane 16 会不会已经把自己改了？"
不会。这是 __shfl_*_sync 这条指令的语义保证的：
参与的 lane 先在这条指令上汇合，每个 lane 读到的都是伙伴在这条指令之前的 sum，再各自做加法写回。
等号右边读到的都是这一轮开始前的旧值，
所以不存在"谁先谁后"。也因此不需要额外的 __syncwarp()：
__shfl_*_sync 自带 warp 内汇合，mask 就是为此而设。
这是 shuffle 比 shared memory 快的根本原因：
shared memory 那条路是"我写内存 → __syncthreads() → 你读内存"，shuffle 是寄存器直连。
先看清执行到第 19 行时还剩什么。
warp_s[0] = warp 0 那 32 个 lane 的和 ⋮ // BLOCK=256 → nWarps = 8 warp_s[7] = warp 7 那 32 个 lane 的和 // 而每个线程寄存器里的 sum 此刻是垃圾：上一段跑完只有各 warp 的 lane 0 是对的， // 那个值已经在第 16 行存进 shared 了。 // 所以第三段只剩一句话要做：把 shared 里这 8 个数归约成 1 个。
① 第 19 行 if (warp == 0)：只剩 8 个数，一个 warp 的 32 个 lane 足够；
其余 7 个 warp（224 个线程）到这里直接退出，不再占用执行资源。它们所在 block 的 shared memory 和 block 槽位，要等整个 block 结束才释放。
warp == 0 等价于 t < 32，
分支边界正好落在 warp 边界上。所以第 22 行那个 0xffffffffu 的承诺成立
（§7·四）；而且分支发生在 warp 之间，warp 内没有线程走两条路，
不产生发散。
"留一个 warp 收尾，其他 warp 退出"是通用模式，
切法必须对齐 warp 边界。
② 第 20 行 sum = (lane < nWarps) ? warp_s[lane] : 0.0f
：注意是 =，不是 +=。
寄存器里原来的 sum 是垃圾，必须整个覆盖；
当成"继续累加"就全错了。这一行同时做三件事：
（一）装载 lane 0…7 各读一格 warp_s[lane]，落在 8 个不同 bank，无 bank conflict。 （二）填单位元 lane 8…31 填 0.0f。mask 承诺了 32 个 lane 到场，得让它们到场又不干扰结果 —— 0 是加法单位元，加几个都不改答案。用单位元把队伍填齐，比精确控制谁参与更简单也更安全。 （三）防越界 warp_s 只分配了 nWarps = 8 个 float（32 字节）。 写成 sum = warp_s[lane]，lane 8…31 会去读 warp_s[8] … warp_s[31]，全部越界。 所以 lane < nWarps 不只是为了填单位元，它还是这段唯一的边界守卫。
③ 第 21–22 行和第 13–14 行逐字相同
，只是这时 32 个 lane 里只有 8 个装着真数据。
设 warp_s[w] = w+1（1…8，总和 36），跟踪 lane 0：
轮 off lane 0 的搭档 搭档的值 lane 0 变成 起始 — — — 1 1 16 lane 16 0 1 + 0 = 1 ← 空转 2 8 lane 8 0 1 + 0 = 1 ← 空转 3 4 lane 4 5 1 + 5 = 6 4 2 lane 2 3+7 = 10 6 + 10 = 16 5 1 lane 1 (2+6)+(4+8) = 20 16 + 20 = 36 ✓ // 5 轮里只有后 3 轮（log₂8）在干实事。想省掉就得按 nWarps 动态算起始 off， // 那要引入运行时计算或分支；而这里的代价只是每个 block 多 2 条 shuffle，还只有一个 warp 在跑。
④ 第 23 行 if (lane == 0) atomicAdd(output, sum);
5 轮走完，只有 lane 0 手里是整个 block 的和，其余 31 个是中间值。
这个 if 确实是 warp 内发散，但只有一条指令，而且 atomicAdd 打在同一地址上本来就串行。
这里 lane == 0 和 t == 0 等价，前者更贴合"warp 坐标系"的语境。
把 N = 2²⁸ 从下往上摊开：同步手段完全由作用域决定
层 1 每线程串行吃 COARSE_FACTOR*2 = 8 个元素 寄存器里的 sum，零同步 8 → 1 层 2 warp 内 32 个 lane → 1 5 轮 shuffle，零同步 256 → 1 层 3 block 内 8 个 warp → 1 shared + __syncthreads() ×1 2048 → 1 层 4 grid 内 131072 个 block → 1 atomicAdd 131072 → 1 验算：131072 block × 2048 元素 = 2²⁸ ✓
warp 内用 shuffle（寄存器直连），零同步；
跨 warp 时 shuffle 够不着，只能用 shared memory + __syncthreads()；跨 block 连 __syncthreads() 都没有，只剩原子加。
R6 改的是 block 内部的第 2、3 层，第 4 层一个字没改：
原子加次数 = grid 大小，R5 和 R6 都是 131072 次。
R4 是 524288 次，是它们的 4 倍。这是 COARSE_FACTOR 的功劳，
和 shuffle 无关。R6 省的是 block 内的 __syncthreads() 和 shared memory 流量。
为什么必须是"三段"，以及一个刚好封闭的边界
第一段 8 个 warp 并行，各自 32 → 1 全在寄存器，0 次屏障 ↓ 8 个数必须过一趟 shared ← 卡点只在这里 第三段 1 个 warp 把 8 → 1 全在寄存器，0 次屏障 // 根本原因一条：__shfl_* 的作用域是 warp，跨 warp 交换它无能为力。 // 所以中间那次 写/屏障/读 省不掉 —— 这就是屏障只能从 8 次降到 1 次、而不是 0 次的原因。
边界刚好封闭：BLOCK 最大 1024 → nWarps 最大 32 → 正好填满一个 warp。 如果 block 能开到 2048 个线程，nWarps 就是 64，一个 warp 装不下， 第三段还要再分一次，结构会变成四段甚至递归。 CUDA"每个 block 最多 1024 个线程"这条硬限制，正好让三段式封闭。
⚠ 第 17 行的 __syncthreads() 必须放在
if (warp == 0) 外面。
它要求 block 内全部 256 个线程都到达；挪进 if，
另外 7 个 warp 永远到不了，结果是死锁。反过来看：正因为它在外面，
第 20 行的 lane 0 才能安全读到 warp 7 写的 warp_s[7]。
R5 树里的那一行和 R6 的这一行是同一棵树，只是"格子"换了地方。
| R5 shared memory 树 | R6 warp shuffle | |
|---|---|---|
| 我的累加器 | s[t] shared memory 第 t 格 | sum 我自己的寄存器 |
| 我要加的搭档 | s[t + stride] | __shfl_down_sync(…, sum, off) lane t+off 的寄存器 |
| 每轮折半的距离 | stride = 128, 64, …, 1 | off = 16, 8, 4, 2, 1 |
| 挡住多余线程 | if (t < stride) 必须有 | 没有 if —— 理由见下框 |
| 同步 | 每轮一次 __syncthreads() | 不需要 lane 天然同步 |
| 一个 block 的开销 | shared memory 访问 1022 次、__syncthreads() 8 次、占 1 KB | shared memory 访问 16 次、__syncthreads() 1 次、占 32 B |
| 发散的轮数 | 后 5 轮 | 0 轮 |
| 搬的字节 | 4N | 4N 一分没变 |
off 从 16 开始而不是 32，因为第一轮要把 32 个 lane 折成 16 对
——这和 R5 的 stride = blockDim.x/2 同构。
循环正好 5 轮，因为 32 = 2⁵。
为什么 R6 敢去掉 if，R5 不敢：这是两者唯一的实质区别
R5 如果去掉 if (t < stride)，本该退出的线程算出的垃圾会
写进 shared memory；下一轮别人读这些格子就读到垃圾，污染会扩散。
R6 去掉 if 完全无害，因为 shuffle
不写任何共享存储——
lane 16…31 的垃圾只留在自己的寄存器里；它们摆出去给别人读的那份，
是上一轮结束时的合法值。
一句话：R5 用 if 挡住多余线程，
是为了保护 shared memory；R6 不需要挡，因为没有共享的东西会被弄脏。
收益是发散轮数 5 → 0：32 个 lane 从头到尾步调一致，
不发生分支。代价是一半的加法白算了，但那些 lane 本来也闲着。
用真数字验证：设 lane i 的初值 = i+1（即 1, 2, … 32），总和应为 528。跟踪 lane 0：
轮 off lane 0 的搭档 搭档此刻的值 lane 0 变成 此时 lane 0 覆盖了哪些原始 lane
起始 — — — 1 {0}
1 16 lane 16 17 1 + 17 = 18 {0, 16}
2 8 lane 8 34 18 + 34 = 52 {0,16,8,24}
3 4 lane 4 68 52 + 68 = 120 上面 4 个，各自 +4 → 8 个
4 2 lane 2 136 120 +136 = 256 16 个
5 1 lane 1 272 256 +272 = 528 全部 32 个 ✓
// 每一轮 lane 0 的搭档，都是"另外半棵树的根"—— 它手里的值恰好是那半棵树的全部和。
2026-09-21 实跑，RTX 4090 (sm_89) · 128 SM · 峰值带宽 1008 GB/s · 每 SM 1536 线程。
六个 kernel 全部校验通过，
compute-sanitizer --tool racecheck 报告 0 hazards。
编译期检查也没有问题：全是 0 bytes stack frame、0 spill，寄存器最多 20 个。
| kernel | 原始耗时 | 减去空 kernel | 相对 R1（净） |
|---|---|---|---|
| 空 kernel | 0.0018 ms | — | 基线 |
| R1 朴素 | 0.0040 ms | 2.2 µs | 1.00× |
| R2 收敛 | 0.0023 ms | 0.5 µs | 4.40× |
| R3 shared | 0.0025 ms | 0.7 µs | 3.14× |
这张表只能看定性结论
净耗时只有 0.5–2.2 µs，
而空 kernel 本身就要 1.8 µs。
减掉空 kernel 后，剩下的量级和 cudaEvent 的分辨率同阶，
±0.1 µs 的抖动就能让倍数翻一番。按原始耗时算，R1→R2 只快 1.74×，
R2→R3 是 0.92×（R3 略慢）。
两种口径的定性结论相同：
R1→R2 变快了，R2→R3 没有变快。
| kernel | 耗时 | 有效带宽 | 达成率 | 相对上一行 | 原子加 |
|---|---|---|---|---|---|
| 下限 4N ÷ 1008 GB/s | 1.0652 ms | 1008 GB/s | 100% | — | — |
| R4 分段 | 1.1203 ms | 958.5 GB/s | 95.1% | 距下限 +5.2% | 524,288 |
| R5 粗化 | 1.1211 ms | 957.8 GB/s | 95.0% | 0.999× | 131,072 |
| R6 shuffle | 1.1212 ms | 957.7 GB/s | 95.0% | 1.000× | 131,072 |
| R5 @ BLOCK=1024 | 1.1224 ms | 956.7 GB/s | 94.9% | 0.999× | 32,768 |
R4 一步就到 95.1%，后面三行一个百分点都没再动。 1.1203 / 1.1211 / 1.1212 ms 三个数的差别在噪声范围内。
| # | 预测 | 标准 | 实测 | 判决 |
|---|---|---|---|---|
| ① | R1→R2 有明显加速 | 净 > 2× | 净 4.40×，原始口径 1.74× | 押中 数字脆弱 |
| ② | R2→R3 再快一截 | 净 > 1.5× | 0.71× —— R3 反而慢 | 没押中 |
| ③ | R3→R4 数量级跳变 | 每元素 > 50× | 292×（净口径也有 82×） | 押中 远超标准 |
| ④ | R4→R5 收益很小 | < 1.2× | 0.999× | 押中 不是"小"，是零 |
| ⑤ | R5→R6 收益更小 | < 1.1× | 1.000× | 押中 同样是零 |
第 ③ 条是核心，它押中得毫无悬念。 注意：这一条必须按每元素耗时比。 实验 B 的 N 是实验 A 的 13 万倍，直接比总时间没有意义。 第 ④⑤ 条预测"收益很小"，实测的收益不是"小"， 是零。
"没收益"可以有三种原因，得排除掉后两种才算预测成立： (a) 带宽真的满了，上面没空间。 ✓ 这就是答案。 R4 达成率 95.1%，而 1008 GB/s 是理论峰值 —— R4 已经站在墙上了。 (b) 原子加本来就不是瓶颈，所以减少它没用。 ✓ 顺带也证实了。 原子加次数 524,288 → 131,072 → 32,768，跨度 16 倍， 耗时 1.1203 / 1.1211 / 1.1224 ms —— 纹丝不动。 (c) 测量噪声把收益盖住了。 ✗ 排除。 实验 B 在毫秒量级、跑了 11 次取平均，三者最多相差 2.1 µs（0.19%）。
所以第 ④⑤ 条不只是押中了，而且是按预测时给的理由押中的。 对账要核对的正是这一点。
// N = 2048 时，全部数据只有 8 KB。 // 第一次读完就整个躺在 L1 里了 —— R2 那些"反复读全局"，实际上一次 DRAM 都没碰。 // 于是 R3 省掉的那些全局访问，本来就是免费的。 // 而 R3 额外付出了：shared 的写 + 读、以及一次额外的 __syncthreads 协调。 // 净效果：持平到略负。
写第 ② 条预测时，我犯了 §1·一 警告过的错：拿书的指标当时间。 R3 的价值在 N=2048 时体现不出来。N 大到装不进 cache 时它才有用， 而那时已经必须用 R4 的分段结构了。
方法论教训：第 ① 条的 4.40× 也不可靠
净耗时只有 0.5–2.2 µs，空 kernel 本身就要 1.8 µs，
减掉之后的量级和计时分辨率同阶。
按原始耗时算只快 1.74×，没达到 2× 的标准。
定性结论（R1→R2 变快了）成立，具体倍数不成立。
下次做这类实验，
要把 N 加大到启动开销可以忽略，或者用 ncu 直接读 kernel 内部计数器。
BLOCK = 1024 （每 SM 只放得下 1 个 block，1024/1536 = 66.7%） 1.1224 ms
BLOCK = 256 （100%） 1.1211 ms
差 0.1%
这条结论有前提，而且方向和 SGEMM 相反
只有当占用率不足导致延迟藏不住时，它才是问题。 归约是纯顺序读，66.7% 的 warp 已经能喂饱显存带宽， 再补上剩下的 33% 也没有带宽可用。 报告 01 §7 的 K5 把占用率从 75% 降到 33.3%， 反而快 1.66 倍。那边是"降低占用率换寄存器复用"， 这边是"提高占用率什么也换不到"。方向相反，道理相同：看瓶颈在哪。
开跑前的担心可以划掉："524,288 次原子加全打在同一个地址上，会不会堵死？"
现代 GPU 的 atomicAdd 在 L2 里完成，
524,288 次分摊到 1.1 ms 里，耗时上看不出影响。
别误读 ptxas 的 "used 1 barriers"
它的意思是占用了 1 个 barrier 硬件资源，不是"执行了 1 次 __syncthreads()"。
R5 循环里的几次 __syncthreads() 共用同一个 barrier 资源，
所以 R5 和 R6 的这一行完全相同。
ptxas -v 看不出
__syncthreads() 次数的减少。
SGEMM 是推理里最费时间的算子，归约是最频繁的算子。 Transformer 的每一层、每一个 token 都要跑好几次归约。
| 它出现在哪 | 是什么形式的归约 | 这一章的哪一课用得上 |
|---|---|---|
| LayerNorm / RMSNorm | 每一行求和 + 求平方和 | 一个 block 管一行，就是 R3 的结构；行长几千，一个 block 正好装得下 |
| Softmax | 两次归约：先求 max，再求 exp 的和 | max 归约可复现，sum 归约不可复现 |
| FlashAttention | 在线 softmax：边扫 KV 分块边维护 running max / sum | 这一章的直接延伸 —— 本质是"能增量合并的归约" |
| 张量并行 allreduce | 跨 GPU 求和（NCCL ring / tree） | 同一棵树，换了一层硬件：lane → warp → block → GPU → 节点 |
| 采样 / top-k / top-p | max、排序、前缀和 | 归约是前缀和的基础 |
| 量化的 scale | 每个 tensor / 每一行求 absmax | 把加法换成 max 的归约 |
FlashAttention 靠的是可合并的归约
FlashAttention 能不落盘地算 attention，靠的是 "softmax 的归约可以分块做、再合并"这个性质。 能不能合并，取决于算子是否满足结合律。 §3·二 那张图下面提醒过这一点。
报告 01/02（SGEMM） 算术强度 0.250 → 12.720，靠分块/粗化/向量化把复用做出来 本报告（归约） 算术强度 0.250 恒定，一步也挪不动 // 同一个优化，在两边的命运完全相反： 线程粗化 SGEMM K3→K4 3.359× 归约 R4→R5 0.999× 提高占用率 SGEMM K5 降到 33.3% 反而快 归约 66.7% 和 100% 打平
SGEMM 教你怎么把复用做出来，归约教你在没有复用可做时该看什么。 推理里一大半算子属于后者，它们都卡在带宽上。 所以"算一遍下限、量一次达成率"这个动作， 比任何单个优化技巧都值钱。
nvcc -O3 -arch=sm_89 -lineinfo reduce_ch10.cu -o reduce_ch10 # 实验 A：单 block，N = 2048（先拿空 kernel 作启动开销基线） ./reduce_ch10 none 2048 1000 ./reduce_ch10 r1 2048 1000 ./reduce_ch10 r2 2048 1000 ./reduce_ch10 r3 2048 1000 # 实验 B：整卡，N = 2²⁸ ./reduce_ch10 r4 268435456 11 ./reduce_ch10 r5 268435456 11 ./reduce_ch10 r6 268435456 11 ./reduce_ch10 r5 268435456 11 1024 # occupancy 对照 # 编译期体检 + 竞态检查 nvcc -O3 -arch=sm_89 -Xptxas -v -c reduce_ch10.cu 2>&1 | grep -A2 reduce_r compute-sanitizer --tool racecheck ./reduce_ch10 r6 2048 1
三个踩过的坑
① R1、R2 都原地修改 input，跑完输入就毁了。
多次计时要每轮重新拷贝输入。
② __syncthreads() 必须写在 if 外面。
写在里面，部分线程永远等不到，就会死锁。
③ unsigned 的 stride 不要写成减到 0 还会再减的形式。
无符号数会绕回最大值，变成死循环。
原始数据一条没删：
Results/Reduction/reduce_20260921.txt 是当天的原样终端输出。
本报告里每个数字都由构建脚本从它现场解析，没有一个是手抄的。
换一张卡，这里所有绝对数字都会变；
但"R4 之后达成率不再动"这个定性结论，在任何带宽受限的卡上都应该成立。
两项还没测
这一批没有测两项：小 N（2¹⁶）下 R5/R6 的对比，以及
cuobjdump -sass | grep SHFL 的条数。
所以"R6 编出了 SHFL 指令"目前只是推断。
报告 01 §8 当初对 float4 也是这样，那边后来用反汇编证实了，
这边还没有。