Kestrel_VLLM - 24 · GDN 线性注意力上 GPU：递推状态、conv 历史与往返成本

24 · GDN 线性注意力上 GPU：递推状态、conv 历史与往返成本

一句话：Gated DeltaNet（GDN）是线性注意力，它每一步都要读一份递推状态、算一次更新再写回，还随行一段卷积历史——这是本系列第一次遇到「有状态递推算子」；本篇记录它上 GPU 的全过程：先用 f32 把七步算子序列钉死，再把三个大投影换成 raw-q4，补上逐算子回读，最后用同内核 A/B 推翻前文一句关于往返成本的定性结论。 前置：建议先读 第 23 篇（raw-q4 内核与常驻池）。 环境：x86-64 + NVIDIA RTX 5090；模型 Qwen3.5-27B（48 层 GDN + 16 层 full attention）；导出 gdn_l0 oracle（f32 权重 + 真实前向）。 三个数：f32 档对 oracle 的 worst max_rel = 3.800e-07（T=5、判据 1e-4；第 六 节「f32 档实测」）；同内核 A/B 量出每层往返 ≈0.01–0.04 ms（48 层 ≈0.5–2 ms/token）⇒ 推翻前文 48×F≈13 ms 的定性结论（第 十一 节）；q4 单层内核三次迭代 0.268 → 0.224 → 0.156 ms（累计 1.72×）、折算有效带宽 238 → 410 GB/s（第 十二 节）。 怎么读：结论看［一、问题与结论］、［十一、往返成本 A/B］与［十五、小结］（必读）；机制看［三/四、核心机制］与［十二、按形状拆开 gemv］（想深挖再看）；［六、f32 档实测］［八、q4 档］［十、逐算子回读］的实测输出块与［十三、试过并回退］是过程与原始数据（可跳，逐算子原始读数见文末附录）。

第 23 篇把 raw-q4 dense 内核与独立常驻池落地：权重终于能像 MoE 那样 以 4-bit 原始布局常驻显存，不再被 int8 展开撑爆。但那条路走的都是无状态算子—— 一次 GEMM 出一份结果，读多少写多少，谁先谁后不影响答案的对错。

27B 的 48 层里还有另一种东西：Gated DeltaNet（GDN）线性注意力。它每一步都要读一份 递推状态 S、算一次 δ 更新、再写回去，还随行一段 conv 历史。这是本系列第一次遇到 「有状态递推算子」——并行怎么切、对拍用什么判据、状态镜像怎么命名与清零， 全都要从零想一遍，不能照抄前四个阶段的无状态 GEMM / attention。

状态：✅ 已落地。

一、问题与结论

二、背景：为什么先做 f32 档，再做 q4 档

gdn_l0 那份 oracle 导出是 f32 权重 + 真实前向（transformers）。这就是先后顺序的全部理由：

• 若这里直接上 q4，误差里会同时混进两样东西：内核写错（量级 O(1)）与 权重量化（rel ~1e-2）。两者混在一起，无法区分。
• f32 档把设备侧算子序列单独钉死之后，第二步的 q4 档只需看差值是否落在量化误差预算内即可。

这与本系列「先隔离再集成」是同一条纪律：一次只引入一个变量， 让每一次失败都有唯一的解释。

本档全部 f32 权重；q4 档（三个大投影走 rq 池）放在 f32 档判据通过之后。

三、核心机制（一）：CPU ↔ 内核逐算子对应

CPU 参考在 src/model/vllm_q35_gdn.c，该文件已由仓库的 GDN 探针验证。 设备内核在 src/npu/cuda/vllm_cuda_kernels.cu。下表把七步（含 ① 里的 β/g 子步）逐列抄全， 「关键点」一列是逐字摘录：

#mermaid-svg-FQXldKhdQCOvsCZG{font-family:\”trebuchet ms\”,verdana,arial,sans-serif;font-size:16px;fill:#333;}@keyframes edge-animation-frame{from{stroke-dashoffset:0;}}@keyframes dash{to{stroke-dashoffset:0;}}#mermaid-svg-FQXldKhdQCOvsCZG .edge-animation-slow{stroke-dasharray:9,5!important;stroke-dashoffset:900;animation:dash 50s linear infinite;stroke-linecap:round;}#mermaid-svg-FQXldKhdQCOvsCZG .edge-animation-fast{stroke-dasharray:9,5!important;stroke-dashoffset:900;animation:dash 20s linear infinite;stroke-linecap:round;}#mermaid-svg-FQXldKhdQCOvsCZG .error-icon{fill:#552222;}#mermaid-svg-FQXldKhdQCOvsCZG .error-text{fill:#552222;stroke:#552222;}#mermaid-svg-FQXldKhdQCOvsCZG .edge-thickness-normal{stroke-width:1px;}#mermaid-svg-FQXldKhdQCOvsCZG .edge-thickness-thick{stroke-width:3.5px;}#mermaid-svg-FQXldKhdQCOvsCZG .edge-pattern-solid{stroke-dasharray:0;}#mermaid-svg-FQXldKhdQCOvsCZG .edge-thickness-invisible{stroke-width:0;fill:none;}#mermaid-svg-FQXldKhdQCOvsCZG .edge-pattern-dashed{stroke-dasharray:3;}#mermaid-svg-FQXldKhdQCOvsCZG .edge-pattern-dotted{stroke-dasharray:2;}#mermaid-svg-FQXldKhdQCOvsCZG .marker{fill:#333333;stroke:#333333;}#mermaid-svg-FQXldKhdQCOvsCZG .marker.cross{stroke:#333333;}#mermaid-svg-FQXldKhdQCOvsCZG svg{font-family:\”trebuchet ms\”,verdana,arial,sans-serif;font-size:16px;}#mermaid-svg-FQXldKhdQCOvsCZG p{margin:0;}#mermaid-svg-FQXldKhdQCOvsCZG .label{font-family:\”trebuchet ms\”,verdana,arial,sans-serif;color:#333;}#mermaid-svg-FQXldKhdQCOvsCZG .cluster-label text{fill:#333;}#mermaid-svg-FQXldKhdQCOvsCZG .cluster-label span{color:#333;}#mermaid-svg-FQXldKhdQCOvsCZG .cluster-label span p{background-color:transparent;}#mermaid-svg-FQXldKhdQCOvsCZG .label text,#mermaid-svg-FQXldKhdQCOvsCZG span{fill:#333;color:#333;}#mermaid-svg-FQXldKhdQCOvsCZG .node rect,#mermaid-svg-FQXldKhdQCOvsCZG .node circle,#mermaid-svg-FQXldKhdQCOvsCZG .node ellipse,#mermaid-svg-FQXldKhdQCOvsCZG .node polygon,#mermaid-svg-FQXldKhdQCOvsCZG .node path{fill:#ECECFF;stroke:#9370DB;stroke-width:1px;}#mermaid-svg-FQXldKhdQCOvsCZG .rough-node .label text,#mermaid-svg-FQXldKhdQCOvsCZG .node .label text,#mermaid-svg-FQXldKhdQCOvsCZG .image-shape .label,#mermaid-svg-FQXldKhdQCOvsCZG .icon-shape .label{text-anchor:middle;}#mermaid-svg-FQXldKhdQCOvsCZG .node .katex path{fill:#000;stroke:#000;stroke-width:1px;}#mermaid-svg-FQXldKhdQCOvsCZG .rough-node .label,#mermaid-svg-FQXldKhdQCOvsCZG .node .label,#mermaid-svg-FQXldKhdQCOvsCZG .image-shape .label,#mermaid-svg-FQXldKhdQCOvsCZG .icon-shape .label{text-align:center;}#mermaid-svg-FQXldKhdQCOvsCZG .node.clickable{cursor:pointer;}#mermaid-svg-FQXldKhdQCOvsCZG .root .anchor path{fill:#333333!important;stroke-width:0;stroke:#333333;}#mermaid-svg-FQXldKhdQCOvsCZG .arrowheadPath{fill:#333333;}#mermaid-svg-FQXldKhdQCOvsCZG .edgePath .path{stroke:#333333;stroke-width:2.0px;}#mermaid-svg-FQXldKhdQCOvsCZG .flowchart-link{stroke:#333333;fill:none;}#mermaid-svg-FQXldKhdQCOvsCZG .edgeLabel{background-color:rgba(232,232,232, 0.8);text-align:center;}#mermaid-svg-FQXldKhdQCOvsCZG .edgeLabel p{background-color:rgba(232,232,232, 0.8);}#mermaid-svg-FQXldKhdQCOvsCZG .edgeLabel rect{opacity:0.5;background-color:rgba(232,232,232, 0.8);fill:rgba(232,232,232, 0.8);}#mermaid-svg-FQXldKhdQCOvsCZG .labelBkg{background-color:rgba(232, 232, 232, 0.5);}#mermaid-svg-FQXldKhdQCOvsCZG .cluster rect{fill:#ffffde;stroke:#aaaa33;stroke-width:1px;}#mermaid-svg-FQXldKhdQCOvsCZG .cluster text{fill:#333;}#mermaid-svg-FQXldKhdQCOvsCZG .cluster span{color:#333;}#mermaid-svg-FQXldKhdQCOvsCZG div.mermaidTooltip{position:absolute;text-align:center;max-width:200px;padding:2px;font-family:\”trebuchet ms\”,verdana,arial,sans-serif;font-size:12px;background:hsl(80, 100%, 96.2745098039%);border:1px solid #aaaa33;border-radius:2px;pointer-events:none;z-index:100;}#mermaid-svg-FQXldKhdQCOvsCZG .flowchartTitleText{text-anchor:middle;font-size:18px;fill:#333;}#mermaid-svg-FQXldKhdQCOvsCZG rect.text{fill:none;stroke-width:0;}#mermaid-svg-FQXldKhdQCOvsCZG .icon-shape,#mermaid-svg-FQXldKhdQCOvsCZG .image-shape{background-color:rgba(232,232,232, 0.8);text-align:center;}#mermaid-svg-FQXldKhdQCOvsCZG .icon-shape p,#mermaid-svg-FQXldKhdQCOvsCZG .image-shape p{background-color:rgba(232,232,232, 0.8);padding:2px;}#mermaid-svg-FQXldKhdQCOvsCZG .icon-shape .label rect,#mermaid-svg-FQXldKhdQCOvsCZG .image-shape .label rect{opacity:0.5;background-color:rgba(232,232,232, 0.8);fill:rgba(232,232,232, 0.8);}#mermaid-svg-FQXldKhdQCOvsCZG .label-icon{display:inline-block;height:1em;overflow:visible;vertical-align:-0.125em;}#mermaid-svg-FQXldKhdQCOvsCZG .node .label-icon path{fill:currentColor;stroke:revert;stroke-width:revert;}#mermaid-svg-FQXldKhdQCOvsCZG :root{–mermaid-font-family:\”trebuchet ms\”,verdana,arial,sans-serif;}

读写

读写

① 四投影 GEMV ×4

β/g：b→β, a→门控

② conv+SiLU

③ 切分 q/k/v

④ q/k 归一

⑤ delta 递推

⑥ gated-RMSNorm

⑦ out_proj

conv 历史 ≈5.9 MB

状态镜像 S ≈151 MB

图 1：GDN 上 GPU 的七步递推链与两处状态。设备内核按序对应 CPU 的 q35_gdn_step：① 四个投影（★ in_proj_b 出 β、in_proj_a 出门控 a，与直觉相反）→ β/g（★ 门控在对数域 −exp(A_log)·softplus(a+dt_bias)）→ ② depthwise 因果卷积 + SiLU（★ 窗口 [s0,s1,s2,x_t] 最旧在前）→ ③ 切成 q/k/v 视图 → ④ q/k 归一（★ 先按 k-head 归一）→ ⑤ delta 递推（α 缩放 → S·k → (v−·)·β → 秩一加 → 用更新后的 S·q）→ ⑥ gated-RMSNorm → ⑦ out_proj。虚箭头是两处需跨 token 保持的设备状态：递推状态 S（[nl][vh][kd][vd] = 48 层 × 3.15 MB ≈ 151 MB）与 conv 历史（[nl][ck-1][conv_dim] ≈ 5.9 MB），合计 ≈157 MB。⑤递推与②卷积无跨线程归约、逐字同序，差异只来自 ①⑦/④/⑥ 三处归约树，故判据取容差而非位级（f32 档 worst max_rel 3.800e-07）。

有三个点是反直觉的，值得单独拎出来（写错任何一处，误差都会是 O(1)）：

这两处「反直觉」在源码里就是两个小核，值得原样贴出来对照：

/* 文件名：src/npu/cuda/vllm_cuda_kernels.cu —— β/g：b 出 β、a 进门控，且 g 在对数域 */
__global__ void k_gdn_beta_g(const float *__restrict a,
const float *__restrict b,
const float *__restrict a_log,
const float *__restrict dt_bias,
float *__restrict beta, float *__restrict g,
int vh) {
const int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i >= vh) return;
const float xa = a[i] + dt_bias[i];
float sp;
if (xa > 20.0f) sp = xa; /* 防 exp 溢出，同 CPU 的分支 */
else if (xa < –20.0f) sp = expf(xa);
else sp = log1pf(expf(xa));
g[i] = –expf(a_log[i]) * sp; /* ★ 对数域（负值） */
beta[i] = 1.0f / (1.0f + expf(–b[i])); /* ★ β 来自 b */
}

/* ② depthwise 因果卷积 + SiLU，并在同核内推进历史（历史存 conv 的输入，非输出） */
__global__ void k_gdn_conv_silu(const float *__restrict raw_in,
float *__restrict out,
float *__restrict hist,
const float *__restrict conv_w,
int conv_dim, int conv_k) {
const int ch = blockIdx.x * blockDim.x + threadIdx.x;
if (ch >= conv_dim) return;
const int nprev = conv_k – 1;
const float *cw = conv_w + (size_t)ch * (size_t)conv_k;
const float raw = raw_in[ch]; /* conv 的输入（历史要存的就是它） */
float acc = 0.0f;
for (int j = 0; j < nprev; j++)
acc += cw[j] * hist[(size_t)j * conv_dim + ch];
acc += cw[nprev] * raw;
/* 历史推进：先读完再写（同线程内顺序保证），左移一帧、末帧放 raw */
for (int j = 0; j + 1 < nprev; j++)
hist[(size_t)j * conv_dim + ch] = hist[(size_t)(j + 1) * conv_dim + ch];
if (nprev > 0) hist[(size_t)(nprev – 1) * conv_dim + ch] = raw;
out[ch] = acc / (1.0f + expf(–acc)); /* silu(acc) */
}

四、核心机制（二）：并行划分与「差异来自哪里」（订正过一次）

并行划分决定了判据形态：

• ⑤ 递推：一个 block 一个 v-head，**thread = dv（独占 S 的一列）**⇒ 对 dk 的三遍扫描 全在线程内完成，既不需跨线程归约、也不会争同一地址；
• ② conv：每线程独占一个通道；
• ④ q/k 归一：一个 warp 一个 k-head；
• ⑥ gated-RMSNorm：一个 block 一个 v-head。

★ 差异来自哪里（这一条订正过一次，别记错）：

⑤ 递推（含 δ 的秩一更新）与 ② 卷积在这里没有任何跨线程归约，逐元素算式与 CPU 逐字相同，单看它们是位级可复现的；真正变序的是 ①⑦ 投影（warp 树形归约）、 ④ q/k 归一、⑥ gated-RMSNorm（shared 二分） 这三处的归约树。 不要把差异归给 dk 求和——它是串行的。

这段订正必须原样保留，因为它直接决定了判据：

• 变序的地方只有三处归约树（①⑦ / ④ / ⑥），它们重排浮点加法，重排就换位；
• 而递推与卷积是串行的、逐字相同的，理论上单看它们可以位级复现；
• 因此整体判据是容差而非位级（与仓库的 GDN 探针同一判据）。

换句话说：GDN 的对拍判据不是「逐字节相同」，而是「相对误差落在归约树能解释的量级内」。 这也解释了为什么 f32 档的 3.8e-7 是个好结果——只差三处归约树就该是这个量级； 哪一步写错，量级会跳到 O(1)。

⑤ 递推的内核就是这条「串行、逐字可复现」纪律的落点：每个线程独占 S 的一列， 对 dk 做三遍串行扫描，顺序严格照 CPU：

/* 文件名：src/npu/cuda/gdn_delta_twin.inc（引擎与探针共享的唯一一份源）
* 一个 block 一个 v-head，thread = dv 独占 S 的第 dv 列 ⇒ dk 三遍扫描全在线程内。 */
__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;
if (vi >= vh) return;
extern __shared__ float sh[]; /* [kd] 存 q 然后 [kd] 存 k */
float *sq = sh, *sk = sh + kd;
const int rep = vh / kh;
const int hh = vi / rep;
for (int d = threadIdx.x; d < kd; d += blockDim.x) {
sq[d] = q[(size_t)hh * (size_t)kd + d];
sk[d] = k[(size_t)hh * (size_t)kd + d];
}
__syncthreads();
float *Sm = S + (size_t)vi * (size_t)kd * (size_t)vd; /* [kd][vd] */
const float alpha = expf(g[vi]);
const float bt = beta[vi];
const float vv = v[(size_t)vi * (size_t)vd + dv];
float kv = 0.0f;
for (int dk = 0; dk < kd; dk++) { /* ① α 缩放 + (S·k) */
float *sp = Sm + (size_t)dk * (size_t)vd + dv;
const float t = (*sp) * alpha; *sp = t; kv += t * sk[dk];
}
const float tmp = (vv – kv) * bt;
float oa = 0.0f;
for (int dk = 0; dk < kd; dk++) { /* ② 秩一加 + (S·q)，用**更新后**的 S */
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;
}

五、状态镜像：S 151 MB + conv 历史 5.9 MB

有状态算子的第一件工程事，是把状态放在哪、多大、怎么清零。

本档是单份镜像（单序列），只提供 gdn_reset（新序列复用）； 投机解码的分支克隆留到后面接 spec_verify 时再加（KV 镜像已有 8 槽池可参照）。

这里有个必须在接线前解决、但本篇只需登记的语义缺口： gdn_step 是 hit-only（未预载就返回 -1，让调用方回退 CPU），但 GDN 的 S / conv_hist 只存在于显存，本轮没有状态导出 API。对无状态算子（GEMM）「miss → 回退」是安全的； 对有状态递推，中途回退会让 CPU 拿着不一致的状态算出错误结果。接线时必须三选一： 保证序列开始前 GDN 权重全常驻（miss 不可能发生） / 加状态导出-导入 / 整序列 all-GPU 或 all-CPU。 这一条与「单份镜像不支持并发与投机分支」一起，是两个准入门槛。

六、f32 档实测

gdn_l0 导出，T=5 token，真权重，–nl 48（下列均为 x86 + RTX 5090 实测）：

[CUDA] gdn: nl=48 dim=5120 k=16×128 v=48×128 ck=4; state 149.6 MB, scratch 0.13 MB
[gdn] tok max_abs_err max_abs_ref max_rel verdict
[gdn] 0 1.907e-06 2.042e+01 9.340e-08 OK
[gdn] 1 9.537e-07 1.251e+01 7.621e-08 OK
[gdn] 2 1.907e-06 1.969e+01 9.688e-08 OK
[gdn] 3 9.537e-07 1.455e+01 6.557e-08 OK
[gdn] 4 5.722e-06 1.506e+01 3.800e-07 OK
[gdn] PASS (0/5 tokens out of tolerance), worst max_rel=3.800e-07
[gdn] reset 后首 token 位级复现（同进程/同发射配置）：不等元素 = 0 / 5120 [相同]
[gdn] 跨层切片（层 0 vs 层 47，同权重、各 5 token）：逐位不等元素 = 0 / 25600 [逐位相同]
[gdn] step(host-facing): 0.400 ms/token/层 有效吞吐 295.5 GMAC/s
[gdn] 本档每 token 读 f32 权重 442.0 MB ⇒ 反算带宽 **1159 GB/s 为上界**（分母含 H2D/D2H/同步，非纯权重读时间；不得当带宽结论引用）

逐条读：

• 正确性（实测，对 oracle）：worst max_rel 3.8e-7，比判据（1e-4）宽 260 倍。 与「只差 ①④⑥ 的归约树」的预期一致；CPU 参考探针同期是 7.8e-7，同量级。
• 性能（实测）：0.400 ms/token/层（另一次运行为 0.450 ms；单次 20 迭代、未做重复次数与 方差统计，只可作量级参考）。反算出的 1159 GB/s 是上界（分母含每层一次 H2D/D2H 往返 + 同步 + 5 次 launch，不是纯权重读时间，不得作带宽 / roofline 结论）。
• f32 档可用的定性结论只有一个：它在显存上根本放不下—— 48 层 × 442 MB = 21 GB/token 的权重读（与第 23 篇的硬阻塞同源）⇒ 必须走 q4。

七、reset 复现与跨层切片：它们是必要条件，不是证明

这两个输出块容易被误读，必须把它们的边界写在前面。

reset 复现：把状态推进 5 个 token 后 reset，重放首 token 与第一次的首 token 位模式相同（0 / 5120，同进程、同发射配置）。★ 这是「清零 + 确定性」的 必要条件检验，不是设备路径的位级证明：

它看不到 S / hist 本身，也不覆盖跨进程 / 跨线程数 / 跨 device。

也就是说，它证明的是「重放一致」，不是「状态结构正确」。必要条件不能当充分条件用。

★ 跨层状态切片（本轮补上的覆盖）：把同一份权重装到第 0 层与第 47 层， 各自 reset 后跑同样的 5 个 token，两串输出 25600 个元素逐位相同（0 / 25600）。

为什么这一条重要——因为它此前根本不存在：

• 切片直接覆盖了 layer*vh*kd*vd 与 layer*(ck-1)*conv_dim 两个偏移量；
• 而在这之前跑的都是 nl=1，layer 恒为 0 ⇒ 偏移恒为 0，等于没测；
• 这与「bump 分配的偏移寻址」是同类风险，而那次是真踩过的。

所以跨层切片不是「多测一条」，而是把「状态偏移寻址」这个此前完全没被覆盖的面补上； 它同时顺带验证了「未预载层被拒绝」（hit-only 生效）。

八、q4 档：三个大投影改走 rq raw-q4 池

落地要点：

• vc_gdn_w_t 新增 rq_in_qkv / rq_in_z / rq_out_w（0 = 该档不启用）； gdn_step 里每张大投影二选一：有 rq key 走 vc_rq_launch_gemv，否则走 k_f32_gemv。
• 新入口 vllm_cuda_gdn_use_rq(layer, key_qkv, key_z, key_out)：全有或全无—— 三条 key 必须都在池里且形状精确匹配，否则整体拒绝、不改状态 （不留「一半 q4 一半 f32」的含糊配置）。
• vllm_cuda_gdn_preload_f32 放宽：三张大投影的指针可为 NULL（表示"本档不供这张"）， 6 张小张量仍必须提供。⇒ q4 档下那 3 个 f32 大缓冲根本不分配（省 442 MB/层）。
• ★ 绝不静默退回 f32：声明了 rq key 就以 rq 为准；池里缺失 / 形状不符 → gdn_step 返回失败，而不是拿一个不存在的缓冲。这与「未预载 → CPU 回退」的语义一致。
• ★ 两条 both 状态由代码强制拦住（注释里的不变量不能只是注释）： preload_f32 在「该投影已切 rq」时拒绝（放在分配之前，避免白分配 442 MB 再拒）； use_rq 在「已有 f32 大投影」时拒绝。否则会出现 both 状态：step 按 rq 走， 而那块 f32 缓冲静默泄漏 442 MB/层。
• ★ 池重建 ⇒ 引用失效：rq_init / rq_clear 会调 vc_gdn_clear_rq_keys， 把各层登记的 key 全部清零。原因：vc_rq_expect 只校验形状；若同一 key 在新池里 绑到另一张形状相同的张量，光靠形状校验会静默用错权重。 （比加 generation 计数器更简单、也更难绕过。）

单投影 q4 误差预算（q35_fused_check –dump 从 model_q4.vqf 导出打包字节； 下面的对照里只差权重量化这一项，不掺别的变量）：

三项合计 61.9 MB（f32 同层是 442.0 MB，7.1×）。

实测（同一次运行里跑两档：层 0 = f32，层 1 = q4；同权重、同输入、各 5 token）：

[gdn] tok max_abs_err max_abs_ref max_rel verdict ← 层 0（f32）
[gdn] PASS (0/5 tokens out of tolerance), worst max_rel=3.800e-07
[gdn] q4 档（第 1 层，三张大投影走 rq raw-q4）：
[gdn] tok=0 vs oracle: maxae=5.556e-01 maxrel=2.721e-02
[gdn] tok=1 vs oracle: maxae=6.421e-02 maxrel=5.131e-03
[gdn] tok=2 vs oracle: maxae=1.296e-01 maxrel=6.581e-03
[gdn] tok=3 vs oracle: maxae=5.014e-02 maxrel=3.447e-03
[gdn] tok=4 vs oracle: maxae=3.923e-02 maxrel=2.605e-03
[gdn] q4 vs f32 档（同权重、只差量化）：maxae=5.556e-01 maxrel=2.721e-02
[gdn] q4 档量级闸门（maxrel < 5e-2）：OK
[gdn] 层 0 (f32): 0.467 ms/token 有效吞吐 253.3 GMAC/s
[gdn] 每 token 读权重 442.0 MB ⇒ 反算带宽 993 GB/s 为上界
[gdn] 层 1 (q4): 0.300 ms/token 有效吞吐 394.0 GMAC/s
[gdn] 每 token 读权重 63.9 MB ⇒ 反算带宽 223 GB/s 为上界

一条口径声明（必须写明）：本档 GPU 走的是 q4 权重 + f32 激活（k_rq_q4_gemv_f32act）。 而 CPU 的生产路径 q35_gdn_step → q35_wv_gemv 在 q4 分支里会先把激活量化成 q8 再算。⇒ GPU-q4 与 CPU-q4 不是同一口径（GPU 这条更准，少一层激活量化误差）。 因此：① 这一节的数字不能与 CPU q4 路径的数字直接比； ② 要与出货口径对齐必须另做 GPU-q4 vs CPU-q4 的对拍。

九、正确性证据是什么（订正：原表述是个恒等式）

这一节是本篇最该记住的一段——因为最初写下的"证据"其实是一个恒等式，什么都证明不了。

原表述（错）：q4 vs f32 档 与 q4 vs oracle 两行数值完全相同 ⇒ 差全部来自量化 ⇒ 连线正确。

为什么是空的：因为 f32 档对 oracle 只差 8e-6 量级，对任意 rec1 都有 |rec1−rec0| ≈ |rec1−exp|（差 ≤1.6e-5），两行必然相等——接错时它也会相等。

真正能证明连线正确的是另外两件事：

量级闸门：q4 档 maxrel 2.72e-02 < 1e-1。★ 这个阈值只是 O(1) 探测器： 连线接错的量级是 1e0（比它大一个数量级），而实测合法值 2.72e-02 距它 3.7×， 留余量以免对正确实现误报 FAIL。它不是"精度合格"判据。

★ 未决项（当时不要当成已通过）：tok0 的 maxrel 2.72e-02 是最坏单投影 （out_proj 6.998e-03）的 3.9×；tok1…4（5.1e-3 / 6.6e-3 / 3.4e-3 / 2.6e-3） 则贴着单投影预算。注意 tok0 的绝对误差 0.556 也是 tok1（0.064）的 ~9×， 不是分母假象。当时的合理解释是「tok0 递推状态为 0、gated-RMSNorm 的归一化因子最小 ⇒ 量化误差被归一化放大」，但本轮没有验证——要定论必须做逐算子 dump 对拍。 在补齐之前，q4 档只能算「连线正确、量级未越界」，不能说"精度满足要求"； 而且 48 层的 2.7e-2 最坏值会层间累积，端到端精度必须另行过闸。

这个未决项，后来在下一节被解除了——不是"放大被解释掉了"，而是 “预算标定的口径本身就错了”。见下节。

十、逐算子回读：解除未决项

为什么先做这个再进下一步：下一步会把 GDN 的驱动方式整个换掉（每层 host 往返 → 设备内常驻）。在那之前若没有逐算子观测，任何数值回归都只能表现为「最后一个数不对」， 无法定位。

为让它可观测而做的改动（数值零变化）：两处原地算子改成出参独立—— k_gdn_conv_silu（保留 ① 原始投影）与 k_gdn_gate_norm（保留 ⑤ 递推输出）， 新增 d_gdn_xcv / d_gdn_onn 两个缓冲。这只是目的指针变了，算式与顺序一字未改； 证据：f32 档端到端仍是 worst max_rel 3.800e-07（与上一档逐位相同的判据值）。

新增入口 vllm_cuda_gdn_get(layer, what, dst, cap)，what 覆盖 {XC,XCV,Z,Q,K,A,B,BETA,G,ON,ONN,OUT,S,HIST}；dst == NULL 为元素个数查询模式； 装不下则拒绝而不是截断。不引入任何每步开销（按需 cudaMemcpy）。

实测表（gdn_l0，T=5；列 = 该算子自身归一的相对误差； 层 0 = f32，层 1 = q4；maxrel@t0 / rmsrel@t0 / rmsrel@tT-1）：

逐算子原始读数见文末 附录 · 表 A1。

结论（三条，第一条是自洽校验，第二条解开了那个未决项）：

未决项的处置：那个未决项已解除——不是"放大未解释"，而是"预算标定口径错了"。 q4 档的准确定性描述是：整层 RMS-相对 ~3–6e-2，误差源是权重 int4 量化本身， 在本层内部不放大。端到端（48 层串联 + 全模型）是否可接受仍需另行过闸。

十一、往返成本 A/B：推翻前文的一句定性结论

动机：前文只能靠"两点模型"猜固定开销，而那个模型自证不成立 （反解出的带宽超过硬件峰值）。于是先做纯粹的 A/B——它决定「设备内常驻解码值不值得做」这个方向性问题。

做法（保证同内核）：把步进体抽成 vc_gdn_step_dev(s, layer, d_h, d_out, stream) （设备指针、不同步），host-facing 与设备内两条路径共用同一份，因此 kernel 与 每 token 的 launch 数完全相同，唯一差别是往返次数（T 次 vs 1 次）。 新增 vllm_cuda_gdn_run_dev(layer, T, h_seq, out_seq)（1 次 H2D + 1 次 D2H + 1 次同步）。

实测（gdn_l0，T=5，–iters 50）：

层 0 等价性：host-facing(T 次往返) vs 设备内(1 次往返) 逐位不等元素 = 0 / 25600 [逐位相同]
层 1 等价性：host-facing(T 次往返) vs 设备内(1 次往返) 逐位不等元素 = 0 / 25600 [逐位相同]
层 0 (f32)：host-facing 0.420 → 设备内 0.404 ms/token
层 1 (q4) ：host-facing 0.260 → 设备内 0.224 ms/token

• 等价性（实测）：两条路径逐位相同（0/25600，两档各一次）。这是最强的等价证据—— 在同内核前提下，把往返从 T 次降到 1 次不改变任何一位。
• ★ 往返成本 ≈ 0.01–0.04 ms/层（两次运行给出 0.012/0.012 与 0.016/0.036； 同一流程里的两次差值是噪声，说明这个量级已到 clock() 的分辨极限）。 ⇒ 48 层全程 host-facing 的往返开销是 ≈0.5–2 ms/token， 不是前文那个 48×F≈13 ms。那句话是错的，已就地标注更正。

★ 真正的瓶颈（这才是本节的结论）：把往返扣掉后，q4 档每 token 仍要 0.224 ms 去读 63.9 MB 权重 ⇒ 有效带宽只有 ~285 GB/s，是 5090 峰值（1792 GB/s）的 16%。 不是带宽 bound，也不是往返 bound，而是 kernel 自身。

由这个发现直接驱动的改动（k_rq_q4_gemv_f32act）：

改后实测：

正确性未退化（实测）：maxrel 2.293e-06 → 2.372e-06、RMSrel 1.227e-06 → 1.208e-06 （都还在 1e-6 量级，变化来自求和序由 1 链变 4 链）；q35_rq_check 两条 entry 仍 0/17408 逐位相同；f32 档 GDN 端到端仍是 3.800e-07；q4 档 O(1) 闸门仍 OK； 逐算子表（上节）在打印精度内完全不变（量化误差 1e-2 量级淹没了 1e-6 级的序差）。

对下一步方向的影响：

• 「为了去掉每层往返而做设备内常驻」的理由被削掉了（往返只值 0.5–2 ms/token）；
• 但设备内常驻仍然要做，理由换了：全模型串联时 hidden 每层上下一次 PCIe 仍是浪费 （≈2 ms/token 量级），且后续的异构 dispatch 本来就需要一个设备内的层循环； 只是它现在不是性能关键路径的第一优先；
• 新的第一优先是 raw-q4 内核的访问策略。

十二、按形状拆开 gemv：定位到软件 f16→f32

第一步拆分（同内核、按形状分别测 gemv 纯 kernel 时间）：

三个 gemv 合计 0.151 ms，占 GDN 单层 0.224 ms 的 67% ⇒ 瓶颈确定在这三个 gemv，而非那 8 个非 gemv 内核，也不是往返。

根因：rq_coef 里每个 lane、每个列块都要跑一次软件 d_f16_to_f32 （含 while (m < 0x400u) 的 subnormal 归一化循环）。每 lane 160 个列块 ⇒ 最多 160 次带分支的软件转换，而它摊到的 FMA 只有 1 次。这才是主开销。

改法：换成硬件 __half2float（1 条指令）。两者对 f16→f32 在非 NaN 输入上 精确且唯一，故应逐位等价。软件版保留在源码里作为「硬件指令必须与之相等」的 参考语义（可肉眼核对，而不是口头断言）。

/* 文件名：src/npu/cuda/q4_gemv_twin.inc
* 软件版 d_f16_to_f32 保留为「参考语义」（subnormal 用 113 – shift，不是 127-15-shift）；
* 热路径改用硬件 __half2float。两者对非 NaN 输入精确且唯一 ⇒ 逐位等价。 */
static __device__ __forceinline__ float d_f16_to_f32(uint16_t h) {
uint32_t sign = (uint32_t)(h >> 15) & 1u;
uint32_t exp = (uint32_t)(h >> 10) & 0x1Fu;
uint32_t mant = (uint32_t)h & 0x3FFu;
uint32_t u;
if (exp == 0) {
if (mant == 0) u = sign << 31;
else {
uint32_t m = mant; int shift = 0;
while (m < 0x400u) { m <<= 1; shift++; }
u = (sign << 31) | ((uint32_t)(113 – shift) << 23) | ((m & 0x3FFu) << 13);
}
} else if (exp == 0x1F) {
u = (sign << 31) | (0xFFu << 23) | (mant << 13);
} else {
u = (sign << 31) | ((exp + 112) << 23) | (mant << 13);
}
float f; memcpy(&f, &u, 4);
return f;
}

/* 热路径取值：标度每 lane 自读（warp 广播），转换走硬件 __half2float */
__device__ __forceinline__ float rq_coef(const uint8_t *blk, int m,
int byte_off, int hi) {
__half h; memcpy(&h, blk + m * 2, 2);
const float s = __half2float(h);
const uint8_t byte = blk[byte_off];
int n4 = hi ? (int)(byte >> 4) : (int)(byte & 0x0F);
n4 = (n4 >= 8) ? n4 – 16 : n4; /* int4 two's complement */
return (float)n4 * s;
}

改后实测（同一天、同机、同 dump）：

★ 精度逐位不变（这就是等价性的实测证据）：

• maxrel 2.372e-06 / RMSrel 1.208e-06——与用软件转换的上一次运行完全相同； 三张 GDN 矩阵的 maxrel/RMSrel 也逐项相同（2.193e-06/1.201e-06、2.802e-06/1.171e-06、 1.798e-06/1.361e-06）。
• GDN f32 档端到端仍 3.800e-07（f32 路径不经过 rq_coef，本就应不变）； q4 档 O(1) 闸门 OK；host-facing vs 设备内两条等价性检查仍 0/25600 逐位相同。
• 唯一未覆盖：NaN 输入下硬件指令可能改变 payload（软件版保留符号与尾数）； 实际块标度不会是 NaN，故不影响，但这是等价性的边界，写在这里备查。

往返成本再次确认很小：本轮的 A/B 给出 0.004 ms/层（比上一轮的 0.012 更小， 同样受 clock() 分辨限制）。⇒ 结论不变：往返不是瓶颈。

仍然剩下的差距：964 GB/s 是 5090 峰值的 54%；GDN 三矩阵加权只有 ~780 GB/s。 剩下的部分要靠换访问策略（当前每 lane 每块只拿 1 个字节，且同组 4 行重复读同一 tile ⇒ L1 流量是 DRAM 流量的 4 倍），而不是继续微调循环。

十三、试过并回退：一个 warp 处理 4 行（负结果，勿重做）

动机：按上面的分析，当前每列块要发 3 条 LDG（x / 系数 / 标度）才换 1 次乘加， 且 x 被同一 tile 组的 4 行各读一遍。于是让一个 warp 直接处理 4 行： x 每块只加载 1 次供 4 行共用，每 128 个乘加的 LDG 由 12 条降到 9 条。 每行仍保持原来的 4 链累加顺序 ⇒ 数值预期与上一版逐位相同。

实测（精度确实逐位不变，但性能退步）：

精度三项（maxrel / RMSrel）逐项与上一版完全相同 ⇒ 改动本身没写错。

原因：warp 总数降到 1/4。out_proj 只剩 5120/4 = 1280 个 warp， 而 5090 有 170 SM × 最多 64 warp = 上万 warp 槽位 ⇒ 占用率掉到 ~12%， 并行度不足以隐藏访存延迟，LDG 减少的收益被完全吃掉（行数越少的矩阵越惨）。 ⇒ 这一刀的方向（记账是对的）但必须与「把每行的 warp 数补回来」配套， 否则不能单独做。

处置：已回退到「warp/行 + 4 链 + 每 lane 自读标度 + 硬件 f16→f32」这一版 （并按规矩在源码注释里写明「不要重做这一刀」）。 回退后四项数字恢复到 1713.4 / 1394.9 / 1244.9 / 1459.7 GMAC/s，精度不变。

★ 教训（写给下一次）：在这个 kernel 上，占用率优先于 LDG 计数。 要减少每 warp 的 LDG，正解是保持 warp 数不变、在 warp 内部把 tile 读取向量化（例如 8 个 lane 各取一个 uint32 的系数块再 shuffle 分发）， 而不是把多个行塞进一个 warp。

十四、边界与已知限制

已知边界（f32 档与 q4 档共用，后者另加一条）：

下一步：

CPU 对照（迁移前基线）

• CPU 参考：src/model/vllm_q35_gdn.c:q35_gdn_step（GDN 七步递推）、include/model/vllm_q35_ops.h:q35_rmsnorm_unit / q35_gemv。
• 迁移要点：首次把「有状态递推算子」搬上 GPU——先用 f32 档钉死七步算子序列（β/g→conv+SiLU→切分→q/k 归一→delta 递推→gated-RMSNorm→out_proj）；建状态镜像 S(≈151MB)+conv 历史；递推/卷积串行逐字同序，差异只归给三处归约树。
• 真机验证：命中 E7（q35_gdn_check PASS：5/5 OK、跨层逐位 0/25600、host↔设备位级一致、每层往返≈0.010ms）。

附录：原始数据表

表 A1 · 逐算子回读实测表（q4 vs f32）（原第 十 节）

十五、小结

相关篇目：第 23 篇、第 25 篇 源码与配套资源：本仓库 https://gitee.com/pei-xiaoguang/kestrel-llm-cuda.git； CPU 推理源码 https://gitee.com/pei-xiaoguang/kestrel-llm