26 · 发射空档三步曲：事件计时 → nsys → 判别实验
一句话：上一轮把 raw-q4 GEMV 的访存粒度改成「每 lane 一个 uint32」后，端到端从 27.2 打到 18.9 ms/token（量级 1.44×），却留下一个问题——「为什么是 1.44× 而不是 1.64×」；本篇用三种递进的手段（逐相 CUDA 事件计时 → Nsight 逐核 profile → 不碰引擎的判别实验）去量「kernel 之间 GPU 无事可做的间隙」，结论被推翻并订正了两次。
前置：建议先读 第 25 篇（访存粒度与 1.44×）。
环境：x86-64 + NVIDIA RTX 5090；模型 Qwen3.5-27B q4；工具 nsys / ncu；出厂档 build-x64\vllm_cuda.dll。
三个数：nsys cuda_gpu_trace 逐核起止量出发射空档 2.40 ms/token = 12.3%（第 4.3 节）；ncu 显示 k_gdn_delta 只用 DRAM 3.04% / SM 1.19% / 0.02 waves/SM ⇒ 并行度饥饿、降字节数没用（第 4.5 节）；判别实验给出每次发射的空档 = 主机侧 CPU ≈5 µs 与 设备地板 ≈1.1 µs 取大者，模型实测平均 2.05 µs（第 5.2 节）。
怎么读：结论看［一、问题与结论］与［六、串起来的教训］（必读）；机制看［四、第二步 · Nsight 逐核 profile］的 4.3–4.5（空档量化 + ncu 硬证据，想深挖再看）；［三、第一步 · 逐相 CUDA 事件计时］与各节末尾的「审计结论与订正」是过程与原始数据（可跳，逐轮原始读数见文末附录）。
第 25 篇留下一个必须回答的问题：为什么是 1.44× 而不是 1.64×？ 差额被归给了三条推断 （其中一条直指「逐层发射的空隙」）。
本篇是那条推断的清算：用三种递进的手段去量「kernel 之间 GPU 无事可做的间隙」—— 先逐相 CUDA 事件计时，再 Nsight 逐核 profile，最后写一个不碰引擎的判别实验。 结论被推翻并订正了两次：第一次「空隙只占 0.2%」，第二次「发射空档 12.3%」， 第三次才回答「这个空档到底是什么」。
它是本系列方法论的集中体现：每一步都要问「这个测量在量什么」。
状态：✅ 已落地（并两次订正）。
一、问题与结论
- 三版测量的口径不同，结论也不同，且后一版推翻了前一版：
- 第一版（逐相 CUDA 事件计时）：空隙只占 0.2%（0.037 ms/token）——错在口径太窄。
- 第二版（nsys 逐核 profile）：发射空档 2.40 ms/token = 12.3%——第一版的结论被推翻。
- 第三版（判别实验）：「每次发射 ~2 µs 固定成本」不成立——空档与 grid / block / smem / 参数个数无关，但与内核时长强相关。
- 第一版为什么看不见空档：它量的是「相邻两个 event 标记之间的区间」，
而空档被摊进了两侧的相时间里——gap = 层循环 − Σ4 相 只覆盖 63 个层间边界，
0.037 ms 是这条窄口径的上界，不是全 token 的空档总量。
- 空档的两个机制，取大的那个：主机侧 CPU 发射成本 ≈5 µs/次（短核时是瓶颈）； 设备侧 kernel 间地板 ≈1.1 µs/次（长核时主机早已领先）。模型实测的平均 2.05 µs 是两者混合。
- 最硬的一条：ncu 显示 k_gdn_delta 只用 DRAM 3.04% / SM 1.19% / 0.02 waves/SM
⇒ GDN 非投影是延迟 / 并行度瓶颈，不是带宽瓶颈。降字节数没用。
- graph 不是万能的：它能把「主机侧 5 µs」那一半消掉，但消不掉设备地板那一半。 每节点成本还随 grid 上升（grid=1 → 525 ns，grid=4096 → 2454 ns）。
- 可拿回的时间只是「乐观（上限）外插」：CUDA Graph ≈1.8 ms（9.3%）、 融小核 0.39–0.86 ms。未在模型路径实测，禁止当收益引用。
二、背景：为什么要专门量「发射空档」
第 25 篇那 1.44× 的账，卡在一句「推断」上。上一轮给出的三条解释里，有一条是：
逐层发射的空隙：每层约 10 次发射，64 层的发射间隙不随内核变快而缩小。
这条推断的麻烦在于——它看不见。它没有对应的探针：
- 反算带宽（761 GB/s）是把 16 GB 权重除以总时间得到的black-box 均值， 管你时间花在算上还是等待上，它都算进去；
- 相级 --skip 分解只能给「哪一类算子花得多」，而「发射间隙」不是一个算子；
- 微基准（--sweep）跑的是背靠背的内核，根本没有 GPU 空闲的窗口。
换句话说，「发射空档」是关键路径上的一段负空间：它既不属于任何一个核， 也不属于任何一次往返。要量它，必须换一套能看见时间线的仪器。
本篇就是这条线的三步走——而且三步的结论互相打架：
- 设备侧 cudaEvent 分段计时（逐相）→ 空隙 0.2%；
- Nsight Systems 逐核 profile（逐核起止）→ 空档 12.3%；
- 独立判别实验（stream − graph）→ 空档与什么有关。
三、第一步 · 逐相 CUDA 事件计时
上一轮的三条解释都是推断，且「跨轮基线不可自证」这一点也已被暴露。
第一步用设备侧 cudaEvent 分段计时把 18.9 ms/token 逐相拆开，把推断换成实测。
3.1 探针本身：可复现的构建线
- 新开关：-DVC_DIAG_MDL_EVENTS（编译期）+ VLLM_CUDA_MDL_EVENTS=1（运行期），
与 VC_DIAG_MDL_SKIP 同一套门控理由：出厂 DLL 里这条路径打不开。
- 插桩位置（vcuda_dev_model_step）：每层 5 个标记 ⇒ 4 个区间
（in_norm / 子层+残差 / post_norm / MLP+残差），外加 3 个全局标记
（步开始 / 层循环结束 / final norm+lm_head 结束）。标记只入队，本步末尾一次同步读全部。
这套探针就落在 vcuda_dev_model_step 里，用 6 个累加槽把「4 个相 + 尾段 + 层循环」分开，gap 就是「层循环 − 四相和」：
/* 文件名：src/npu/cuda/vllm_cuda_kernels.cu —— 逐相事件计时的收尾（节选）
* 只在 -DVC_DIAG_MDL_EVENTS 构建里存在；每层 5 个标记 ⇒ 4 个区间 + 3 个全局标记。 */
#if defined(VC_DIAG_MDL_EVENTS)
double ph[4] = {0, 0, 0, 0};
for (int L = 0; L < nl; L++) {
const int base = L * 5;
for (int k = 0; k < 4; k++) { /* in_norm / 子层 / post_norm / MLP */
float d = 0.0f;
if (cudaEventElapsedTime(&d, ev[base + k], ev[base + k + 1]) == cudaSuccess)
ph[k] += (double)d;
}
}
double loop_ms = 0.0, tail_ms = 0.0;
cudaEventElapsedTime(&d, ev[nl*5 + 0], ev[nl*5 + 1]); loop_ms = (double)d; /* 层循环 */
cudaEventElapsedTime(&d, ev[nl*5 + 1], ev[nl*5 + 2]); tail_ms = (double)d; /* final+lmhead */
const double sum = ph[0] + ph[1] + ph[2] + ph[3];
s->mdl_ev_gap_ms += (loop_ms - sum); /* ★ gap = 层循环 − 四相和（只覆盖 63 个层间边界）*/
fprintf(stderr,
"[CUDA] mdl-ev AVG n=%.0f | step=%.3f ms | layer-loop=%.3f "
"(in_norm=%.3f sub=%.3f post_norm=%.3f mlp=%.3f | 4ph-sum=%.3f gap=%.3f %.1f%%) | "
"tail(fnorm+lmhead)=%.3f\n", ...);
#endif
- 口径声明：探针不改输出（同 diag DLL、--ref _q4ref64 在探针开/关下三项 maxrel
在打印精度内同值：2.2536e-06 / 2.0334e-06 / 3.3887e-06，argmax 5/5；
⚠ 这是 5 位有效数字的比对，不是逐位比对），但会扰动计时（每步 64×5+3=323 次 record）。
同构建 A/B 实测：探针开 19.611 ms/token vs 探针关 18.903 ms/token ⇒ 探针本身 0.708 ms（3.7%）。
⇒ 所有绝对值只在「探针开」这一组内部互相比较；与上一轮的 18.9 对比时要记住这 3.7%。
构建（新增 -CudaDefines，诊断档放另一个产物目录，绝不覆盖出厂档）：
powershell -File tools\build\build_x64.ps1 -Cuda -OutDir build-x64 # 出厂
powershell -File tools\build\build_x64.ps1 -Cuda -OutDir build-x64-diag `
-CudaDefines VC_DIAG_MDL_EVENTS=1 # 诊断
顺带解决了上一轮记的那个构建坑：-OutDir 传相对路径时，
MinGW ld 收到的 -o 也是相对路径 ⇒ 不再触发「中文路径 → cannot open output file」，
exe 也能正常链接（本轮 rc=0）。旧写法（省略 -OutDir）会把绝对路径喂给 ld，必然失败。
★ 审计订正（6b/6a）：
- 相对 -OutDir 原先在 Test-Path/New-Item 时按 cwd 解析，
而编译器在 Push-Location $root 之后按 **\(root** 解析 ⇒ 从别的 cwd 调用时会「目录建错地方」； 已把 `Push-Location\)root` 提到 OutDir 解析之前，三处基址统一。
- -CudaDefines 非空却把 -OutDir 指向出厂目录时，脚本直接拒绝（原先只有注释提醒）
—— 否则会产出「名为出厂、实为诊断」的同名 DLL。
3.2 实测（全 64 层，iters=30，n=32 平均；同会话内四组背靠背）
⚠ 口径分层（审计 5a 订正）：下表前三行是 cudaEvent 口径（step 是事件窗口内），
不包含任何 host wall 时间；「探针关」那行是工具 clock() 口径（含 H2D/D2H+同步），
两者不是同一个钟，必须分开读。
探针开的 wall 口径（工具 clock()，同一构建、同一 harness）：
19.611 ms/token；探针关：18.903 ms/token ⇒ 探针自身 0.708 ms。
占分母 19.611 的 3.6%、占 18.903 的 3.75%（文稿取 ≈3.7%，两个分母都写明）。
⚠ step(event)=19.482 与 wall=19.611 不是同一个量（差 0.13 ms ≈ h 的 H2D + 输出 D2H
- 主机侧开销），不可互相代入。
三跑稳定性（gap 是三个数里最稳的一个）：混合档第三次独立运行
step=19.510 / sub=9.459 / post_norm=0.534 / mlp=8.644 / gap=0.037 / tail=0.470
（wall 19.642）⇒ gap 三次都是 0.037–0.039，而 sub 在 9.459–9.566、
post_norm 在 0.403–0.534 之间漂 ⇒ 「空隙≈0」这条结论不依赖抖动，
而占比类的数字要按 ±5% 读。
子层的可加性自检：把 --synthetic 得到的两类各层子层耗时按 48/16 加权
（0.16367×48 + 0.11788×16 = 9.74）对实测混合 9.566 ⇒ 1.8%（上一轮同法为 0.4%）。
⇒ 「按层型可加」在本轮成立（与上一轮被否掉的整体线性可加不是一回事）。
3.3 归因表：18.9 ms/token 花在哪
⚠ 本表是「探针开」口径，且未扣除 3.7% 探针开销（审计 4d 订正）。
★ 审计订正（4c）：上一版写的「合计 = 实测 step（差 0.02%）⇒ 归因被实测验证」是错的， 那是定义闭合，不是独立校验。把两行「非投影」的定义代进求和即可看出：
合计 = P总 + (M子层 − P子层) + M_norm + (M_MLP − P_MLP) + M_gap
= P_lm + (M子层 + M_norm + M_MLP + M_gap) = P_lm + M_loop
投影项在求和里被完全抵消，合计恒等于「实测四相和 + 实测空隙 + 预测的 lm_head」。
它与实测 step 的差只等于 P_lm − M_tail。⇒ 0.02% 只说明「预测的 lm_head 0.454 ≈
实测 tail 0.453」这一件事（0.2% 吻合），不是「整张表被独立验证」。
本表唯一的独立对账就是这一条；除此之外「投影 = 字节×速率」始终是外插。
探针开销的量级影响（4d）：GDN 非投影 4.84 恰恰是「探针开的子层实测 − 无探针的微基准预测」
的残差 ⇒ 探针开销会整块计入它。若 0.708 ms 按四相大小比例摊，子层约吃 0.18 ms
⇒ GDN 非投影 ∈ [4.66, 4.84]（全摊给子层则是 [4.13, 4.84]）。
排名结论不受影响（GDN 非投影仍是第一），但引用数字时要带上这个区间。
⚠ 另外 0.708 ms 落在哪一相并未定位（空隙只有 0.037，却总增 0.708）——
它可能在相内的 device 侧 event 成本，也可能是 host 侧 323 次 record + 每步 256 次
cudaEventElapsedTime 的开销。这一层矛盾本轮没有解释。
（tail 0.453 ms ≈ lm_head 单个投影的预测 0.454 ms ⇒ lm_head 已计入「投影」行，
final norm 本身在噪声里。）
3.4 裁决：上一轮的三条推断
- 推断 #2「逐层发射的空隙」被否定：空隙只有 0.037 ms/token = 0.2%。 ★★ 2026-10-01 二次订正：本条的口径太窄，结论在第二版被推翻。 0.037 ms 只覆盖 63 个层间边界；相内、相与相之间的空档被算进了各相时间里。 nsys 实测：每次发射后有 ~2 µs 空档，合计 2.40 ms/token = 12.3%。 ⇒ 「GPU 从头到尾没有等人」是错的；下面那句「融合不是为了省发射」也要改回来。 ⚠ 审计 4a：这 0.037 是层间边界（步起点→第 0 层、63 个层间、末层→循环尾）的测量，
相位内部（子层里 6–8 个核、MLP 里 5 个核之间）的间隙不在这个数里，
它们被算进了各自的相。所以精确表述是：
「在 63 个层间边界上未观测到可测量的空隙（≤0.6 µs/边界，已到 cudaEventElapsedTime 约 0.5 µs 的分辨极限）⇒ 0.037 是上界」，再结合「每 token 约 640 次发射 vs 单 token
19 ms」⇒ 与「主机持续领先、GPU 无空转」一致（是外推，不是直接观测）。
⇒ 「合并小投影以省发射开销」不是为了省发射；
它的收益只能来自内核自身的冷启动/短核效率。attn k/v 的 233 GB/s 因此
不是发射间隙问题（这条也排除了一个替代解释）。
- 推断 #1：由「实测 + 微基准预测」联合定位（审计 4b：上一版写「被证实并量化」过强）： 非投影部分 7.01 ms/token = 36%（= 实测 step 19.482 − 预测投影 12.47），不随访存粒度变快。 最大一块是 GDN 的非投影，∈[4.66, 4.84] ms（24–25%）：conv + β/a + qk-norm + 递推， 48 层 × ≈0.10–0.101 ms/层。 量级与上一轮用另一套数据反算的 0.096 ms/层一致（+5%） —— 但两者都是 「实测 − 预测」，不是完全独立。
- 推断 #3（attn k/v）保留：它占 0.2 ms 量级，且已被上面第 1 条限定为「核内效率」。
因此 1.44× 与「投影 1.63×」并不矛盾：只有 64% 的时间是会变快的那部分。 按本轮分解算结构性上限：
投影@v4 ≈ 12.47 × 1.6335 = 20.38 ms 其余 = 7.01 ms（不随本改动变快）
上限 = (20.38 + 7.01) / (12.47 + 7.01) = 27.39 / 19.48 = 1.406×
★ 这条上限反过来给上一轮的跨会话数字提供了一个（弱）独立支撑：
27.39 = 实测后总量 + 微基准预测的投影节省量（12.47×0.6335），把它与
归档的改动前 27.2 比，差 0.7% —— 一端是本会话实测、另一端是
上一会话归档，中间靠微基准模型连接 ⇒ 不是纯循环。
⇒ 「基线无法自证」的顾虑被削弱（未消除：会话标识不可回溯、
归档端自身有 27.145/27.214/27.572 的抖动、且这条链仍依赖 1.6335 这个外插值）。
★ 本节不改写上一轮记的 1.20× 与 1.44×： 本轮 wall 口径 18.888/18.903 与上一轮的 18.916/18.935 在 ±0.8% 内 ⇒ 1.44× 不动； 1.20× 本轮未复测，保持原样（仍是单次运行、无方差）。
3.5 对路线图的影响（按可拿回的时间排序）
★ 但这三项都还带一个未辨识项（审计 7，必须前置成落地门禁）：
0.10 ms/层 里究竟是延迟还是吞吐，本轮没有区分。48 个 GDN 层在时间轴上是强串行
（S_t 依赖 S_{t-1}），若这 4.8 ms 主要是小核延迟 / 依赖链而不是带宽，
优化手段会完全不同（减核/融合 vs 提带宽），照「带宽」去改会白干。
⇒ 在动 GDN 非投影之前，先用逐核（nsys 或减记录版探针）把 4.8 ms 再拆一层，
并先做一次不带探针的确认（因为 4.84 同时继承了探针开销与微基准模型误差）。
3.6 审计结论与订正（第一步）
审计只做静态核对（本轮未复跑；它明确声明了这一点），但抓到了一处方法学硬伤。
审计的总结论：改动（探针）达标；结论方向达标，但「被独立验证过」这个更强的主张必须订正—— 「为什么是 1.44× 而非 1.64×」这个核心答案站得住，需要修的只是它的证据成色。
四、第二步 · Nsight 逐核 profile
第 3.5 节定了前置门禁：动 GDN 非投影之前，先用逐核手段判断那 4.8 ms 是延迟还是吞吐。
第二步用 nsys（逐核时间 + GPU 时间线）+ ncu（SpeedOfLight / Occupancy）把它做实——
顺带把第一步的结论推翻了。
4.1 方法与环境坑（可复现）
# ★ nsys 的 protobuf 拒绝非 UTF-8 的 working_directory ⇒ 仓库根含中文时**直接失败**
# （libprotobuf ERROR ... working_directory contains invalid UTF-8 data）。
# 解法：subst 一条全 ASCII 路径，所有参数都走它。
subst X: <repo_root>
nsys profile --stats=true --force-overwrite=true -o X:\nsys_gdn64 `
X:\GitHupSRC\tools\build\out\q35_gpu_model.exe X:\Modl\Qwen3.5-27B-q8local `
--vqf X:\Modl\Qwen3.5-27B-q8local\model_q4.vqf --layers 64 `
--lib X:\GitHupSRC\build-x64\vllm_cuda.dll --iters 1
nsys stats --report cuda_gpu_kern_sum --format table X:\nsys_gdn64.nsys-rep # 逐核汇总
nsys stats --report cuda_gpu_trace --format csv X:\nsys_gdn64.nsys-rep # 逐核起止（算空隙）
ncu --kernel-name regex:k_gdn_delta --launch-count 1 --section SpeedOfLight --section Occupancy <exe> ...
- ⚠ -o X:\... 的产物落在「仓库根」（X: 映射的就是仓库根）：实测 .nsys-rep ≈ 20 MB、 .sqlite ≈ 320 MB。跑完必须删（本轮已删），否则仓库里躺 340 MB。
更省事的写法是把 -o 指到 %TEMP%。
- ⚠ 步数账目（审计 8 订正）：上表命令行没有 --ref ⇒ 实际是 3 个 token 步
（T=1 首算 + 预热 + 1 计时），401×3 = 1203 次 rq gemv ✓ 与下表对得上。
审计复跑时带了 --ref _q4ref64 ⇒ 7 步 / 2807 次 rq，折算后逐项与本节表吻合 ≤1%。
- 被剖析的运行：--layers 64 --iters 1，出厂档 build-x64\vllm_cuda.dll（无探针，两轮都没开 VLLM_CUDA_MDL_EVENTS）。
- ⚠ 这套工具是诊断用，不进任何构建线；它不改引擎一个字节。
4.2 逐核时间表（nsys cuda_gpu_kern_sum，3 步聚合后除以 3）
一条真正的独立对账（这是第 3.3 节缺的那一条）：nsys 实测 q4 投影 12.256 ms/token
vs 微基准 cold 速率外插的 12.47 ⇒ 差 1.7%。
⇒ 「投影 = 字节 × cold 速率」这个模型独立于「定义闭合」被外插对账了
（⚠ 措辞比上一版收窄：预测端仍含合成权重的不确定性）。
k_f32_gemv 的 96 次/token 也精确对上「β/a 各一次 × 48 个 GDN 层」。
这两核的「形状」就是本节要看清的对象——k_gdn_delta 一个 block 一个 v-head、每线程 128 步串行扫描：
/* 文件名：src/npu/cuda/gdn_delta_twin.inc —— 递推内核（ncu 显示 DRAM 3.04% / SM 1.19%）
* grid = 48 block（一个 v-head 一个）、block = 128 thread（thread = dv 独占 S 一列）；
* 下面两条循环各 128 步，是**时间轴上的串行链**。 */
__global__ void k_gdn_delta(float *__restrict S, const float *__restrict q,
const float *__restrict k, const float *__restrict v,
const float *__restrict g, const float *__restrict beta,
float *__restrict o, int kh, int kd, int vh, int vd) {
const int vi = blockIdx.x;
const int dv = threadIdx.x;
...
float kv = 0.0f;
for (int dk = 0; dk < kd; dk++) { /* ① α 缩放 + (S·k)：128 步串行，每步读+写全局 S */
float *sp = Sm + (size_t)dk * (size_t)vd + dv;
const float t = (*sp) * alpha; *sp = t; kv += t * sk[dk];
}
...
for (int dk = 0; dk < kd; dk++) { /* ② 秩一加 + (S·q)：再 128 步 */
float *sp = Sm + (size_t)dk * (size_t)vd + dv;
const float t = *sp + sk[dk] * tmp; *sp = t; oa += t * sq[dk];
}
o[(size_t)vi * (size_t)vd + dv] = oa;
}
β/a 的投影则是另一种极端——只有 6 个 block：
/* 文件名：src/npu/cuda/f32_gemv_twin.inc —— β/a 的 GEMV（rows=48、cols=5120）
* 一个 warp 一行，lane 以 32 为步长扫列；grid 只有 6 个 block（wpb=8）。 */
__global__ void k_f32_gemv(const float *__restrict W, const float *__restrict x,
float *__restrict y, int rows, int cols) {
const int warp = (int)((blockIdx.x * blockDim.x + threadIdx.x) >> 5);
const int lane = (int)(threadIdx.x & 31u);
if (warp >= rows) return;
const float *wr = W + (size_t)warp * (size_t)cols;
float acc = 0.0f;
for (int c = lane; c < cols; c += 32) acc += wr[c] * x[c];
for (int off = 16; off; off >>= 1)
acc += __shfl_down_sync(0xffffffffu, acc, off);
if (lane == 0) y[warp] = acc;
}
4.3 ★ 发射空档：每 token 2.40 ms（12.3%）
Direct 量（同一次运行内，cuda_gpu_trace 的逐核起止）。审计独立复跑复现（同会话、同一 boot）：
逐轮原始读数见文末 附录 · 表 A1。
★ 「这 2.4 ms 不是 CUPTI 造出来的」有两条独立论证（第一条是审计建议补的、更硬）：
- wall − Σ(内核时长) 与追踪工具无关：审计复跑在无 nsys 下测得 wall 18.942 ms/token
（同一会话、同一 exe/DLL），而 nsys 下的 busy 只有 16.4–17.0 ms ⇒ 差 ~2.2 ms。
若这 2.2 ms 是 CUPTI 加的，那么无 nsys 时内核必须比 nsys 下快 13% —— 不成立
（CUPTI 的时长是 GPU 硬件时间戳，追踪开销加在主机发射路径上，不会拉长内核）。
- 跨度 18.96/19.43 与无探针 wall 18.888/18.903 同值；若每次发射被凭空加 2 µs， 跨度应膨胀到 ~21 ms。
⚠ 残留不确定性：两条都建立在「CUPTI 不改变内核时长」之上；彻底闭合需要非 nsys 的 busy 测量 （例如给第一步的探针做「减记录版」）。⇒ 证据强度高、非绝对。
- 空档个数 1168/token，均值 2.03–2.05 µs、中位数 2.02 µs（分布很窄）。
- ⚠ 窗口口径披露：上面的「跨度」定义为「该步第一次 rq gemv（第 0 层的
本文由 AI图鉴自动抓取并翻译，版权归原作者所有；如需查看原文请点下方链接。
• Launching of an empty three-step song: event timing →nsys → separation experiment
One sentence: the last round, after replacing the raw-q4 GEMV access particle content with "one uint 32 per lane", the end-to-end call from 27.2 to 18.9 ms/token (1.44 x volume) left a question - "Why 1.44 x instead of 1.64 x"; this section uses three progressive means (when measured by CUDA events) to measure "the gap of nothing between the GPUs between the kernels", which was overturned and revised twice.
Prefix: It is recommended to read section 25 (visiting particle size and 1.44 x).
Environment: x86-64 + NVIDIA RTX 5090; model Qwen3.5-27B q4; tools nsys / ncu; commission-x64\vllm cuda.dll.
Three figures: nsys cuda gpu race measured launch slots on a nuclear-by-nuclear basis 2.40 ms/token = 12.3% (Section 4.3); ncu shows that k gdn delta used 3.04% / SM 1.19% / 0.02 waves/SM parallel hunger and drop byte numbers only (Section 4.5); separate experiments gave each launch space = CPU ≈5 μs on the host side and ≈1.1 μs on the equipment floor, with an average of 2.05 μs measured in real terms (Section 5.2).
How to read: Conclusions read [I, issues and conclusions] and [VI, lessons learned in combination] (required reading); Mechanism 4.3-4.5 (scrutinized hard evidence of quantization + ncu) read by [IV, II] Nsight read by verification; [III, first step] The "Audience Conclusions and Revisions" at the end of the sections are process and raw data (which can be typified in the original reading round to see the appendix at the end of the text).
Title 25 leaves a question that must be answered: Why 1.44 x instead of 1.64 x? The difference is attributed to three inferences (one of which refers directly to the "scatter-by-screech space").
This is the inference of liquidation: three progressive means to measure "the gap between the kernels and the GPU that can do nothing" -- first by CUDA events, then Nsight by nuclear protocol, and finally by writing a separate experiment without touching the engine. The conclusions were overturned and revised twice: for the first time, "0.2 percent of the space" and for the second, "12.3 percent of the launch slot," and for the third time, "What is this gap?"
It is a central expression of this series of methodological theories: each step asks, "What is this measurement doing?"
Status: ✅ landed (and revised twice).
I. ISSUES AND CONCLUSIONS
- The three versions of the survey varied in calibres and conclusions and the latter reversed the previous version:
- The first edition (CUDA-by-CUDA event time): the gap is only 0.2% (0.037 ms/token) — the error is too narrow.
- Second edition (nsys nuclear-by-nuclear): launch slot 2.40 ms/token = 12.3% — The conclusions of the first edition were overturned.
- The third edition (discretionary experiment): "The fixed cost of ~2 μs per launch" is not valid - the gap is not related to the number of grid / block / smem / parameters, but to the size of the kernel.
- Why can't you see the blanks in the first edition: It's "the area between two events next to each other."
And the blank space is spread across both sides of the spectrum - the gap = layer cycle - the thorium 4 phase covers only 63 layers of boundary.
0.037 ms is the upper limit of this narrow calibre, not the total space of the full token.
- Two mechanisms in blank, the larger one: the CPU launch cost on the side of the mainframe ≈5 μs/team (short-nucle time is a bottleneck); and the ≈1.1 μs/team on the side of the device ≈ernel floor ≈1.1 μs/team (the host is already ahead of the long core). The average of 2.05 μs measured in the model is a mixture of the two.
- Hardest article: ncu shows k gdn delta only 3.04% / SM 1.19% / 0.02 waters/SM
GNN non-projection is a delay / parallel bottleneck, not a bandwidth bottleneck. The drop of byte is useless.
- graph isn't all-powerful: it can get rid of the half of the "host side 5 meters" but not the half of the floor of the device. The cost per node also increases with grid (grid = 1 → 525 ns, grid = 4096 → 2454 ns).
- The time taken back is just "optimistic": CUDA Graph ≈1.8 ms (9.3%), fusion nucleus 0.39-0.86 ms. It is not measured in the model path and is forbidden to be quoted when it comes to proceeds.
II. Background: Why a special "launch slot"
Book 25, 1.44 x's account, stuck in a "presumed". Of the three explanations given in the last round, one is:
Sphere-by-layer space: about 10 launches per layer, and 64 layers of space does not shrink with the speed of internal nuclear transformation.
The problem with this assumption is that it can't be seen. It has no corresponding probe:
- Inverse bandwidth (761 GB/s) is the equal Black-box value obtained by dividing 16 GB weights by the total time, whether you spend time counting or waiting;
- Phase-skip decomposition can only cost "what type of algorithms" and "launch gap" is not a algorithm;
- The micro-baseline (-sweep) runs with a back-to-back kernel, with no windows available for GPU.
In other words, the "launch gap" is a negative space on the critical path: It doesn't belong to any of the cores, nor to any of the round trips. To measure it, an instrument that can see the time line must be replaced.
This is the three steps of the line — and three steps of the conclusion fighting each other:
- Equipment side cudaEvent timed (by phase) 0.2% space;
- Nsight Systems, nuclear-by-nuclear policy (start-to-date) slot 12.3 per cent;
- What's the gap with the test of independence?
CUDA Events Timed
The three interpretations of the previous round were extrapolated and the "cross-cycle baseline is not self-evident" has been exposed.
As a first step, 18.9 ms/token is to be removed from one side of the device by the cudaEvent sub-section and extrapolation is to be replaced by a measurement.
3.1 Probe itself: Recoverable construction Line
- New switches: -DVC DIAG MDL EVENTS (compilation period) + VLLM CUDA MDL EVENTS=1 (operational period),
The same set of door-control reasons as VC DIAG MDL SKIP: This path cannot be opened in the DLL.
- Plug position (vcuda dev model step): 5 markers per layer 4 compartments
(in norm / sublayer + disability / post norm / MLP+ disability), plus 3 global tags
(Step Start / Circle End / final / final norm+lm head). Marks are entered only in line and read all at the end of this step.
This probe falls in the vcuda dev model step, and separates the "4 symmetry + end + layer cycle" with a six-symmetric slot, and gap is the "Cyber cycle - 4 symmetry":
* File name: src/npu/cuda/vllm cuda kernels.cu - End of each event by event (extract)
* Only exists in the -DVC DIAG MDL EVENTS construction; 5 tags per layer ⇒ 4 slots + 3 global tags. */
#if defined(VC_DIAG_MDL_EVENTS)
double ph[4] = {0, 0, 0, 0};
for (int L = 0; L < nl; L++) {
const int base = L * 5;
For (int k = 0; k < 4; k++){/ * in norm/ sublayer/ post norm / MLP*/
float d = 0.0f;
if (cudaEventElapsedTime(&d, ev[base + k], ev[base + k + 1]) == cudaSuccess)
ph[k] += (double)d;
}
}
double loop_ms = 0.0, tail_ms = 0.0;
cudaEventElapsedTime (&d, ev [nl*5 + 0], ev [nl*5 + 1]; loop ms = (doule)d; /* layers */
cudaEventElapsedTime(&d, ev[nl*5 + 1], ev[nl*5 + 2]); tail_ms = (double)d; /* final+lmhead */
const double sum = ph[0] + ph[1] + ph[2] + ph[3];
s->mdl ev gap ms + = (loop ms-sum); /* ggap = layer cycle - quadrilateral sum (covering 63 layers)*/
fprintf(stderr,
"[CUDA] mdl-ev AVG n=%.0f | step=%.3f ms | layer-loop=%.3f "
"(in_norm=%.3f sub=%.3f post_norm=%.3f mlp=%.3f | 4ph-sum=%.3f gap=%.3f %.1f%%) | "
"tail(fnorm+lmhead)=%.3f\n", ...);
#endif
- caliber declaration: do not change the output of probes (same as dig DLL, -ref q4ref64 with three maxrel on/off
The same value within the print accuracy: 2.2536e-06 / 2.0334e-06 / 3.3887e-06, argmax 5/5;
⚠ This is a 5-bit ratio, not a one-by-one comparison), but it can disturb the time of the count (64x5+3 = 323 times per step).
Also constructed A/B fact sheets: Probes 19.611 ms/token vs probes 18.903 ms/token gill probes themselves 0.708 ms (3.7%).
⇒ All absolute values are compared only within the group of `prospecting needles'; this is kept in mind in comparison with the previous round of 18.9.
Construction (new - Cuda Defines, another catalogue of products, never covered by the plant):
Powershell-File tools\buld\bild x64.ps1-Cuda-OutDir built-x64 #
powershell -File tools\build\build_x64.ps1 -Cuda -OutDir build-x64-diag `
-Cuda Defines VC DIAG MDL EVENTS=1 # Diagnosis
When the construction pit that was written in the previous round was solved, the --OutDir passed the relative path.
MinGW ld received -o is also a relative path ⇒ → cannot open output file
Exe also has a normal link (Rc=0 for this round). Old writings (abbreviated-OutDir) will feed the absolute path to ld and will fail.
Audit revision (6b/6a):
- Relative - OutDir used to parse by cwd while Test-Path/New-Item,
The compiler then uses **(root**) to analyze the ⇒ from another cwd call to "write the directory"; `Push-Location\)root ' has been referred to OutDir before the OutDir analysis has been made.
-Cuda Defines refused the script when it was not empty and pointed it out of the factory catalogue.
– Otherwise, it will produce the same DLL as the "worker, actually diagnosed".
3.2 Measured (all 64 floors, sets = 30, n = 32 average; four back-to-back groups in the same session)
⚠ caliber stratification (audit 5a revision): The first three lines of the table below are the cudaEvent caliber (step is within the event window).
It does not contain any host Wall time; the line "probate" is a tool lock() caliber (with H2D/D2H+Sync).
They are not the same clock and must be read separately.
Wall caliber (toollock(), same build, same Harness):
19.611 ms/token; probe level: 18.903 ms/token ⇒ probe itself 0.708 ms.
3.6% of the denominator of 19.611 and 3.75% of the denominator of 18.903 (scripts ≈3.7%, both denominators specified).
⚠step(event)=19.482 and Wall=19.611 are not the same amount (0.13 ms ≈ h) H2D+ output D2H
- Host side expenses) cannot be applied to each other.
Triple running stability (gap is the most stable in three digits): third independent operation of the hybrids
step=19.510 / sub=9.459 / post_norm=0.534 / mlp=8.644 / gap=0.037 / tail=0.470
(wall 19.642) ⇒ lap three times is 0.037-0.039, while sub in 9.459-9.566,
The conclusion that post norm drifts between 0.403 – 0.534 is that the "space zero" does not depend on shaking.
And the percentage is to be read at ±5%.
Additional self-censorship of the sublayer: time-consuming for the two layers of the --synthetic
(0.16367 x 48 + 0.11788 x 16 = 9.74) Mixing of measurements 9.566 ⇒ 1.8% (0.4% for the previous round).
⇒ ⇒ )。 )。 )。 )。 )。 )。 )。 )。 )。 )。 )。 )。 )。 )。 )。 )。 )。 )。 )。 )。 )。 )。 )。 )。 )。 )。 )。 )。 )。 )。 )。 )。 )。 )。 )。 )。 )。 )。 )。 )。 )。 )。 )。 )。 )。 )。 )。
3.3 Attribution table: 18.9 ms/token
⚠ This table is a `screech' caliber, without deduction of 3.7% of the cost of the probe (audit 4d revision).
• Audit Revision (4c): The "total = measured step (0.02% difference) ⇒ attribution is wrong ", as written in the previous version, which is defined as closed, not independent verification. It is clear from the two lines of "non-simulation" that:
Total = P total + (M sublayer - P layer) + M norm + (M MLP - P MLP) + M gap
=P lm+ (M + M norm + M MLP + M gap) =P m + M loop
The projection item is completely offset in sum, and the sum of constants is equal to the "Mem head of four-phase measurements and + gap measurements" predicted.
It's only equal to the difference between measured step and P lm - M tail. ⇒ 0.02% only indicates 'm head 0.454 ≈
This is not an "independent verification of the whole table":
The only independent reconciliation in this table is this one; otherwise, the " Projection = byte Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z Z.
Quantification effect of the cost of probes (4d): GDN non-projection 4.84 is exactly the "substrate reality of probes - micro-base projection without probes"
The disability of the gill probe costs the whole piece of it. If 0.708 ms is distributed at a four-size scale, the submersible is about 0.18 ms
⇒ GNN non-fissile [4.66, 4.84] (the whole amortization layer is [4.13, 4.84]).
The results of the ranking are not affected (GDN non-projection remains the first), but the figures are quoted in this compartment.
⚠ Also 0.708 ms in which phase is unpositioned (0.037 in gap, overall increase 0.708) -
It can cost events on the side of the picture, or it can host 323 record + 256 steps
CudaEvent ElapsedTime's expenses. This layer of contradiction is not explained.
(Tail 0.453 ms ≈m head forecast 0.454 ms ⇒m head is included in the projection line.
(In the noise)
3.4 award: three inferences from the previous round
- Insumption #2 "Step by layer" is denied: only 0.037 ms/token = 0.2%. 2026-10-01 Second revision: this article is too narrow in calibre and its conclusions were overturned in the second edition. 0.037 ms covers only 63 inter-story boundaries; gaps between and within phases are counted over time. nsys: ~2 μs blank after each launch, total 2.40 ms/token = 12.3%. ⇒ "GPU without a man" is wrong; the following sentence, "Integration is not intended to save launch", is also to be replaced. Audit 4a: This 0.037 is a measurement of the inter-story boundary (step starting layer 0, 63 layers, end layer loop),
The gap within the phase (between 6-8 cores in the sublayer and 5 in the MLP) is not in this number.
They're counted into each other. So the exact expression is:
"No measurable gap (0.0.6 μs/boundary, reaching the resolution limit of cudaEvent ElapsedTime approximately 0.5 μs) ⇒0.037 is upper bounds" combined with "approximately 640 launches per token vs single token
19 ms" is consistent with "the host continues to lead and the GPU is empty" (excuse, not direct observation).
⇒;;;;;;;;;;;;;;;;;;;
It can only benefit from its own cold start/short nuclear efficiency. 233 GB/s attn k/v
It is not a question of the launch gap (this article also excludes an alternative explanation).
- Insumption #1: co-positioned by `real + microbase projection' (audit 4b: previous version of `verified and quantization' too strong): Non-projection 7.01 ms/token = 36% (= actual step 19.482 - projected projection 12.47), faster than available. The largest piece is a non-projection of GDN, ∈ (4.66, 4.84] ms (24-25%): conv + beta/a + qk-norm + push, 48 floors x ≈0.10 – 0.101 ms/ layer. The scale is consistent with the previous round of 0.096 ms/layers (+5%) with the other set of data, but both are "real-forecasting" and are not completely independent.
- Infer that #3 (attn k/v) is retained: it represents 0.2 ms and has been qualified by Article 1 above as “nuclear internal efficiency”.
Therefore, 1.44 x does not conflict with " Projection 1.63 x": only 64% of the time is the part that gets faster. Structural ceiling based on current cycle breakdown:
Projection
Upper limit = (20.38 + 7.01) / (12.47 + 7.01) = 27.39 / 19.48 = 1.406 x
This ceiling, in turn, provides a (weak) independent basis for the last round of cross-session numbers:
27.39 = total after measurement + projection savings for microbaselines (12.47 x 0.6335), and
Before the file was changed, 27.2 to 0.7 percent less -- one end of the conversation, one end of the story, one end of the story.
The last conversation was archived, and the middle was connected to a micro-baseline model, and the platinum was not pure cycle.
⇒ The concern that "baseline cannot be self-proven" is weakened (unsolved: session marking is non-retroactive,
The filing end itself has 27.145/27.214/27.572 vibrating and the chain is still dependent on 1.6335, the external plugin value).
• This section does not rephrase the 1.20x and 1.44x: The current round's wal cal. 18.888/18.903 and the previous round's 18.916/18.935 within ±0.8% the ⇒1.44x remain intact;
3.5 Impact on the road map (in order of recoverable time)
But all three also carry an undiscovered item (audit 7 must be pre-empted):
0.10 ms/layer, whether delayed or swallowed, does not distinguish between the current cycle. 48 GDN layers are strong lines on the time axis
If this 4.8 ms is mainly a small nuclear delay / a dependent chain rather than a bandwidth,
Optimizing means can be completely different (reduced/integrated vs bandwidth) and change to "bandwidth" for nothing.
⇒ Before moving the GDN non-film, use a nuclear-by-nucle (nsys or less-recorded probe) to remove 4.8 ms and tear one layer down.
And first, a confirmation without a probe (because 4.84 also inherited the cost of the probe and the miscalculation of the microbase model).
3.6 Audit findings and revisions (first step)
The audit performed only static reconciliations (the current round was not repeated; it clearly stated this), but a methodological wound was caught.
The audit's general conclusion: the change (probation) met the target; the conclusion reached the target, but the stronger proposition of " been independently validated" had to be revised - the core answer, "Why is 1.44 x instead of 1.64 x", stood, and all that needed to be repaired was the colour of its evidence.
Nsight, nuclear-by-nuclear
Section 3.5 provides for a pre-door ban: before moving the GDN non-projection, a nuclear-by-nuclear determination of whether the 4.8 ms is delayed or swallowed.
Step two uses nsys (check-by-check time + GPU time line) + ncu (Speed OfLight / Occupancy) to make it real --
The conclusion of the first step was reversed.
4.1 Methods and environmental pits (repeated)
# Protobuf of nsys refused to work directory ⇒ repository containing Chinese**
# （libprotobuf ERROR ... working_directory contains invalid UTF-8 data）。
# Solve: subst a full ASCII path, all parameters follow it.
subst X: <repo_root>
nsys profile --stats=true --force-overwrite=true -o X:\nsys_gdn64 `
X:\GitHupSRC\tools\build\out\q35_gpu_model.exe X:\Modl\Qwen3.5-27B-q8local `
--vqf X:\Modl\Qwen3.5-27B-q8local\model_q4.vqf --layers 64 `
--lib X:\GitHupSRC\build-x64\vllm_cuda.dll --iters 1
Other Organiser
Other Organiser
ncu --kernel-name regex:k_gdn_delta --launch-count 1 --section SpeedOfLight --section Occupancy <exe> ...
- ⚠-o X:...the product falls on the "storey root" (X: the map is the warehouse root): the results are real.nsys-rep ≈ 20MB,.sqlite ≈ 320 MB. After running must be deleted (this round has been deleted) or 340 MB lies in the warehouse.
It's easier to write by pointing -o to %TEMP%.
- ⚠ Step (Audit 8 Revised): No -ref ⇒ in table command line above
(T=1 initial + preheat + 1 time), 401 x 3 = 1203 rq gemv thalles against the following table.
Ref q4ref64 ⇒7 step / 2807 rqs, converted to line by item 1 per cent of the table in this section.
- Anticipated operation: -layers 64-its 1, out of the factory build-x64\vllm cuda.dll (without a probe, two rounds of VLLM CUDA MDL EVENTS).
- This set of tools is diagnostic and does not enter any construction line; it does not change the engine to a byte.
4.2 Nuclear-by-nuclear schedule (nsys cuda gpu kern sum, 3-step aggregation divided by 3)
A true independent reconciliation (this is the one missing in section 3.3): nsys based q4 projection 12256 ms/token
vs Micro-base cold 12.47 ⇒ 1.7% variance
⇒ Projection = byte x cold speed = model independently of "Definition closed" is plugged in.
(The wording is narrower than the previous version: the prediction end still contains the uncertainty of synthetic weights).
96 times/token of k f32 gemv is also precise in relation to "beta/a each x 48 GDN layers".
The shape of these two cores is the object of this section. -k gdn delta a block a v-head, 128-step scan:
* File name: src/npu/cuda/gdn delta twin.inc - Gradual kernel (ncu display 3.04% / SM 1.1%)
* Grid = 48 block (one v-head), block = 128 track (thread = dv column S);
* The following two cycles, 128 steps each, are the chain on the time axis**. */
__global__ void k_gdn_delta(float *__restrict S, const float *__restrict q,
const float *__restrict k, const float *__restrict v,
const float *__restrict g, const float *__restrict beta,
float *__restrict o, int kh, int kd, int vh, int vd) {
const int vi = blockIdx.x;
const int dv = threadIdx.x;
...
float kv = 0.0f;
For (int dk = 0; dk < kd; dk+){/* 1 α Zoom + (Sk): 128 steps per step + write global S*/
float *sp = Sm + (size_t)dk * (size_t)vd + dv;
const float t = (*sp) * alpha; *sp = t; kv += t * sk[dk];
}
...
For (int dk = 0; dk < kd; dk+){/ * 2 + 1+ (Sq): 128 additional steps*/
float *sp = Sm + (size_t)dk * (size_t)vd + dv;
const float t = *sp + sk[dk] * tmp; *sp = t; oa += t * sq[dk];
}
o[(size_t)vi * (size_t)vd + dv] = oa;
}
The beta/a projection is another extreme - only 6 block:
* File name: src/npu/cuda/f32 gemv twin.inc - beta/a GEMV (rows=48, cols=5120)
* One row of warp, lane sweeps at 32 steps; grid has only 6 blocks (wpb=8). */
__global__ void k_f32_gemv(const float *__restrict W, const float *__restrict x,
float *__restrict y, int rows, int cols) {
const int warp = (int)((blockIdx.x * blockDim.x + threadIdx.x) >> 5);
const int lane = (int)(threadIdx.x & 31u);
if (warp >= rows) return;
const float *wr = W + (size_t)warp * (size_t)cols;
float acc = 0.0f;
for (int c = lane; c < cols; c += 32) acc += wr[c] * x[c];
for (int off = 16; off; off >>= 1)
acc += __shfl_down_sync(0xffffffffu, acc, off);
if (lane == 0) y[warp] = acc;
}
4.3 Firing slot: 2.40 ms (12.3%) per token
Direct quantity (in the same run, cuda gpu race) Audit independent recurrence (same session, same Boot):
See the end of the text. Appendix. Table A1.
★ "This 2.4 ms is not made by CUPTI" has two independent arguments (the first is the harder version of the audit recommendation):
- Wall – thorium (kernel duration) has nothing to do with tracking tools: audit rerun is measured without nsys
(Same session, exe/DLL) and only 16.4-17.0 ms ~2.2 ms under nsys.
If this 2.2 ms is CUPTI plus, no nsys must be 13% faster than nsys -- no.
(CUPTI duration is GPU hardware time stamp, tracking costs added to the host launch path and does not stretch the kernel).
- Range 18.96/19.43 equal to the value of wall 18.888/18.903 without probes; if each launch is emptied with 2 μs, the range should expand to ~21 ms.
Residual uncertainty: both are based on "CUPTI does not change kernel duration"; complete closure requires non-nsys-based measurement (e.g. "reduced version" of the first probe). • High-intensity, non-absolute evidence.
- Vacant 1168 token, average 2.03-2.05 μs, median 2.02 μs (a narrow distribution).
- Window caliber disclosure: "Range" above is defined as "the first rq gemv of the step"
Auto-fetched and translated by AI Atlas. All rights belong to the original author.