第 12 章　负结果、反直觉与未完成项 ​

一条系统路径能走通，不等于所有岔路都通向悬崖。这一章把那些测过之后不成立的尝试、和结果相反直觉的发现、以及当前还没有做完的事收起来——让整份报告有一个可以合上的结尾。

12.1　fp9 kernel 调优里的负结果 ​

整包 prefetch 是负收益 ​

第 10 章已经提过。在所有 expert forward 之前，对整个 expert payload（约 516 KB）做全量 _mm_prefetch(..., _MM_HINT_T0)，在所有 selected_experts 点上都是稳定更慢。expert 数越多，负收益越明显。原因：prefetch 本身额外扫了一遍 payload，占带宽且污染 L3 cache，而 forward 计算本身已经能驱动硬件 prefetch。

scalar-scale 不如向量乘 ​

把 scale 的计算从“内积算完后乘标量”改成“内积前先把 x 的每个分量乘 scale 向量”，在 in_features=64 上反而慢了约 1.3%。原因推测是：提前做 x * scale 需要额外一轮向量乘和一轮 load，而原来的 hsum × scale 只在每 128-row block 结束时做一次标量乘，成本更低。

split2 双 FMA 依赖链是负收益 ​

把 inner dot 手拆成两条独立的 FMA 依赖链（split2），在 in_features=64 上慢了约 5.7%，在 128 上也未见正收益。AVX2 只有 2 条 FMA 单元，理论上拆链能打满，但实际瓶颈不在 FMA 吞吐，而是在 load 和 memory latency。

sign chunk 直读和 aligned prescale 都没有赢 ​

尝试把 K=128 的 sign 读取从 uint64 lo/hi + cmp/cmov/shr 改成按 packed-sign layout 直接按 uint16_t chunk 读；同时把 x_scaled 的预缩放从编译器自动向量化改成显式 _mm256_load_ps/_mm256_mul_ps/_mm256_store_ps。objdump 确认了 codegen 确实变了，但实跑结果：两条路径分别慢了约 0.2% 和 0.6%，合在一起反而更差。

down_proj blocked 不如 row-tiling ​

把 down_proj 64/128 的外层改成 block_row += 128、同一 128-row block 内只构造一次 scale（blocked 模式），在 64 上慢了约 5.7%，在 128 上慢了约 3.5%。不如当前最好的 tile2（相邻多行共享同一份 x load）。

65536-entry full-mask table 是大幅负收益 ​

建一张 65536 项的 __m256i sign-mask 查找表，每个 16-lane chunk 直接按 uint16_t 取完整 ymm 掩码，期望减少 sign 解码的分支和移位。结果在 64 上慢了约 24%，在 128 上慢了约 26%。推测原因是 L1 cache 被这张大表污染，反而把真正需要的权重数据挤出去了。

12.2　系统层路线调查里不成立的路径 ​

daemon 真计算通信重叠尝试先收回 ​

02-prime-2k 跑出 16.171 tok/s 之后，2026-04-08 ~ 2026-04-12 期间又尝试把 dual-lane overlap 往“真计算通信重叠”的 daemon 路线上推进。

设计思路不走 IRQ 和中断驱动，也不直接跳到完整 userspace NIC/protocol stack，第一版选"常驻 decode kernel + core0 net poll + command buffer/deferred response"。JudgeDuck-OS 接了 daemon trap path / step path / bounded netpoll，duck-llm kernel 接了 duck_daemon_entry(...) 和 compute-time poll hook。在 Qwen3-30B-A3B-FP8 tp12 上，daemon poll 一度已经压到每步约 1 次 syscall、单次约 4 µs，局部性能症状确实出现过改善。

但 correctness、lifecycle 和 prepare 后半死状态没有被收住：prepare 看似成功，随后 metadata 1 byte write-buffer 稳定 -233，duck 进入“半死但未彻底掉线”的不可服务状态。多个问题叠在一起——lifecycle/close 语义、host/OS ABI 与统计口径、bring-up 与 stale daemon 清扫、core0 straggler / poll 路径、code 与 notes 认知漂移——最终两个仓库都做了 checkpoint 后回退到 pre-daemon 正确基线。

这轮尝试说明：把 poll 下沉进 compute-time path 这件事本身并非没有收益，理论目标是在 two-batch overlap 条件下让 decode 每 step 再省掉约 20% 时间。但当前这一轮没有把它做成可交付结果。这条线最终只保留为一次有价值但未落地的设计/实现尝试，并未纳入当前正式结果。

V3.2 sparse/indexer 路径回退 ​

2026-03-18 曾短暂接通过一版 indexer + sparse mask 路径，用于完成早期 bring-up / correctness 验证。但实测 decode 约 3.15 tok/s，几乎没有体现出 indexer 该有的加速效果。

原因并不在于 sparse 路线本身“不划算”，而是当时的 active 测试都落在 <=2048 的范围内——在这个区间里 indexer 根本不会被触发，等于白白多了一层 mask 管理和索引构建的逻辑，却看不到任何收益。与此同时，这条额外的路径增加了实现复杂度和潜在的错误面。

所以系统很快回退到更干净的 official_fp8 + <=2048 full-attn 路。2026-03-22 后，旧 legacy / indexer / BF16 kv-cache 实现已从代码入口删除，不再出现在 active 路径中。

attention proj 的专用 FP8 GEMM 路线 ​

早期代码里曾经预留过 attention_proj_mode = "fp8_gemm" 这条专用路径，意图是让 attention 的 q/k/v/o 投影层独立走一块 FP8 GEMM backend。但调查后发现这条路从来没有真正落地——它一直是个占位名字，没有形成可用的 backend。2026-03-20，这组旧参数被显式删除（见第 6 章相关 commit），因为继续保留只会让后续人误以为 attention 还有一条独立的 proj backend 配置面。

实际上，attention 投影层的 FP8 权重保留和计算，已经通过通用 Linear 路径和 DUCK_TRITON_FP8_LINEAR 环境变量在正常工作——不需要一条 attention 专用的 FP8 GEMM 路线。旧占位的删除反而收敛了认知，避免了“两个开关控制同一件事”的混乱。

系统层 FP8 GEMM backend 的调查 ​

要让 attention 和 indexer 投影在 GPU 上走 FP8，需要一个能跑 sm_120（RTX 5090 / Blackwell）的 block-scaled FP8 GEMM backend。这轮调查了几条公开路径，全部在本机环境证伪：

  • torch.nn.functional.scaled_mm：PyTorch 自带的公开接口，对需要的 128×128 block scale 粒度不成立（NotImplementedError: DeepSeek style scaling only supported in SM90），无法覆盖 DeepSeek-V3.2 checkpoint 的 scale layout
  • DeepGEMM：DeepSeek 官方的 FP8 GEMM 库，在本机上卡在 tcgen05.* not supported on sm_120a——Blackwell 架构的 MMA 指令集还没有被 DeepGEMM 支持
  • Transformer Engine（NVIDIA）：在 M=1 的 decode 形状上不适用，无法直接用于逐 token forward
  • trtllm-gen 的 MLA backend：直接报 Unsupported architecture，不支持 sm_120

这四条路径全部走不通之后，系统才转向“自定义 Triton kernel”这条路线。最终 DeepSeek-V3.2 的 attention 投影和 shared experts 走的是通用 Linear 路径 + DUCK_TRITON_FP8_LINEAR=1——不是调用某个现成 FP8 GEMM 库，而是系统自己处理 FP8 权重加载和 Triton compute。

12.3　反直觉发现 ​

expert 数增加时 Gparam/s 反而提升 ​

直觉上 MoE 增加 expert 数会让时间线性增长。实测并非如此：selected_experts 从 4 到 20（参数量 5x），wall time 只涨 4x，Gparam/s 从 13.96 升到 17.30。原因是固定成本（barrier、worker 启动延迟等）被摊薄了。这说明MoE 的瓶颈不是简单的“更多 expert 一定更慢”，而是在有固定开销的前提下，专家数增加反而让 kernel 本体在总时间里占更大比例。

read-only 明显高于 DDR4 理论值 ​

第 10 章详细讨论过：在 N100 + DDR4-3200 单通道上，roofline harness 测到单核 ~17 GB/s、4 核 ~35 GB/s 的 effective read ceiling，高于理论 25.6 GB/s。这个矛盾还没被完全解释清楚，当前口径只能叫“effective logical read ceiling”。

persistent judge 的 AP 晚到等待环 ​

persistent judge 接通后，AP（application processor，即 secondary 核）如果晚到 BSP（bootstrap processor），会把“当前已发布 step”误判成“下一次 step”，导致整个 step 卡死。这个 bug 不是 persistent 模式本身固有的，而是 step 同步语义不够显式（没有区分 active phase 和 idle phase）导致的。修复方式是把 step 同步改成显式的 active/idle 相位，而不是简单加超时。

MTP1 stateless 路径的接受率只有 17% ​

第一版 MTP1 实现按照“MTP layers don't use KV cache”这个注释，走 stateless single-token 路径：attention 只看当前位置，不维护 prefix chain。跑起来后 greedy top-1 接受率只有约 17%（36/208）。根因是：MTP draft 分布和主干 verifier 分布大幅错位。补上 MTP1 prompt prefill + 独立 fp8 KV cache + confirmed-token-only append 后，接受率大幅抬升到约 72%。

non-duck decode 那段时间的去向 ​

在 flattened runtime 之前，MoE decode 约 77ms/token 里，duck 侧 judge_wait 约 47.8ms，剩下的约 29ms 不在 duck 里——其中 attention 的 torch.cat KV cache 追加是最直接可优化的一项。这个发现驱动了 static KV cache → flash attention 2 → flattened runtime 的整条链路。

12.4　roofline 的 caveat ​

为什么不能叫 raw DDR roofline ​

第 10 章的方法学已经排掉了四类可能的测量误差：

  1. 4 核地址切分没算错（disjoint slice）
  2. read loop 没被编译器优化掉（objdump 确认）
  3. libc sanity 量级对得上（read-only > memset > memcpy）
  4. forced-cold + unique windows 后仍然不降

但 35 GB/s（4 核）仍然高于 DDR4-3200 单通道 25.6 GB/s。

可能的解释方向还有：

  • 计时口径问题：当前测的是 core-visible completion time（TSC），不一定等于 memory pins 上真实传输完成时间
  • 平台差异：JudgeDuck / 当前运行环境下的真实可见内存路径，和“单通道 DDR4 raw pin bandwidth”是否完全同一层定义

对当前 duck-llm / MoE fp9 这条线来说，这已经不影响核心判断——真实 MoE fp9 hot 约 20.3 GB/s（wall）只用了 4 核 effective ceiling ~35 GB/s 的大约 58%，主矛盾仍然不在“带宽打满”上。但如果后续有人想用这套 bench 回答“离 25.6 GB/s 还差多少”，需要重新设计更接近 STREAM / IMC 计数器口径的测量。

12.5　当前未完成项 ​

JudgeDuck-OS 独立前史 ​

这份技术报告里的 JudgeDuck-OS 相关章节（第 4 章、第 8 章相关部分）只写了相对于 origin/master 的增量。JudgeDuck-OS 作为独立项目的完整历史、设计初衷、早期评测系统的形态、以及它为什么会有一套完整的 sandbox ELF 执行抽象——这些内容没有写进主文。不是不重要，而是不属于 duck-llm 系统的核心故事。

更系统的图表和 benchmark 覆盖 ​

当前报告主要依赖一次长跑结果（2026-03-29，1024 tokens，质数查询 prompt）和一组后续补跑的 testcase 档案（2026-04-08，7 组 case）。多 prompt decode 性能分布和长输入下的 TTFT / decode 证据已开始成形，但距离标准 benchmark 仍有差距：

  • 缺标准 benchmark 套件的结果（MMLU、HumanEval 等）
  • 缺不同上下文长度下的性能曲线（当前只有两个离散长度点）
  • 缺更细粒度的 per-layer / per-expert 延迟热图

更深的 event-based sync 清理 ​

第 7 章提到 flattened runtime 的 segment 边界仍然需要 torch.cuda.synchronize() 才能安全让 Rust 消费 host buffer。pair overlap runtime 把它改成了 CUDA event，但只在 pair lane 层面。general flattened runtime（非 pair 模式）的 event-based sync 清理还没有最终完成。

更长上下文与 indexer ​

当前 max_model_len = 2048。超过 2048 则需要实现 indexer——这是 DeepSeek-V3.2 的 sparse attention 机制，用于在长上下文下只计算部分 KV 位置。当前 indexer 还没有补，所以 > 2048 的上下文暂时无法在该系统上正确运行。

但即便把 indexer 补上，还有一个更深的硬件层 blocker：在本机 RTX 5090（sm_120）上，Full MLA 的 kernel 支持仍然不完整。2026-03-20 的一轮 probe 确认了几件事：

  • cutlass MLA：当前版本不支持 sm_120
  • trtllm-gen 的 MLA backend：直接报 Unsupported architecture
  • 唯一在 sm_120 上能跑的 xqa kernel：与当前 active 的 contract 不兼容

也就是说，> 2048 目前不只是“差补一个 indexer”的问题——即使 indexer 接通了，底层的 MLA attention backend 在 sm_120 上也还没有可用的实现路径。这是当前系统的硬边界。

更长的上下文还会带来：

  • 更大的 KV cache（显存压力）
  • 更长的 prefill 时间
  • MTP1 的 KV cache 管理复杂度增加

这些都没有系统验证过。

多轮对话的增量 prefill 与会话级 KV cache ​

当前 chat-duck.py 每轮对话都会把完整 chat history 重新 tokenize，再对整段 prompt 做一次 forward。模型层虽然有 prepare_generation_cache() 和 _cached_len 这些基础设施，但 serving 层还没有把“只处理新增输入”做成正式能力。

要支持增量 prefill，至少需要明确三件事：

  1. 会话级 KV cache 的生命周期定义：需要明确一个 session 的 KV cache 何时创建、何时释放、以及放在 GPU / CPU / 磁盘的哪一层。
  2. 增量 prefill 的 dispatch 逻辑：命中缓存后，只对新增 token 做 prefill，而不是整段重算。
  3. 权重缓存与 KV 缓存的边界：当前仓库里已有 duck 的权重缓存（weight cache）和 generation cache 两套机制。多轮对话缓存属于“模型状态缓存”，不是“权重缓存”，文档和实现中必须明确区分，不能混写。

这件事在 TODO 清单里列为 P1。它和缓存设计、serving 协议、模型 state 管理都绑在一起，不是单点 kernel 优化能解决的，而是一个产品级能力的门槛。

fp9 prefill kernel 性能仍待优化 ​

当前的 fp9 格式已经能支撑 duck 的 MLP / MoE / prefill，但 prefill kernel 的性能还远未达到 N100 的理论算力上限。

按 Intel 官方标称频率和 FMA instruction throughput 粗算，N100 单核 FP32 理论峰值约 54.4 GFLOP/s（3.4 GHz × 16 FLOPs/cycle），4 核合计约 217.6 GFLOP/s。而当前实测中，dense prefill 的 matmul 性能约在 12 ~ 44 GFLOP/s 之间——差距显著。

瓶颈主要来自几个方面：

  • 非纯 GEMM 操作占比大：fp9 解码、scale 读取、silu 激活函数、fp32/bf16 类型转换等开销在每次 forward 中都存在，拉低了有效 GFLOP/s。
  • tile scratch 被放在栈上：当前每个 worker 的临时数组放在 1 MiB 的用户态栈上，BATCH_TILE_SIZE 保守固定在 4。如果不把大临时数组挪到共享 runtime scratch，继续放大 tile 有溢出风险。
  • 参数没有做完整 sweep：N_TILE 已经调过（默认收口到 2），但 M_TILE、cache blocking、权重重用策略等还没有系统扫过。

这在最小跑通阶段是合理的优先级选择——最重要的验证目标是 decode 能否跑通，prefill 暂时只追求“可用且不太慢”。但作为下一阶段的优化空间，prefill kernel 的性能提升路径是清晰的：先把 scratch 从栈上移走释放 tile 大小的约束，再系统扫 1/2/4/8 的 tile batch 曲线，同时评估 cache blocking 和预取策略的收益。

项目声明 / Project Disclaimer

本项目为作者以个人身份、利用业余时间推进的个人娱乐项目；除非另有明确说明，它与作者的任何雇主、客户、学校、单位或其他组织均无合作、雇佣、委托、赞助或背书关系，也不代表任何该等主体的立场。除普通个人捐赠外，本项目未获得任何资金支持。

This project is a personal hobby project developed by the author in a personal capacity and in personal time. Unless explicitly stated otherwise, it has no collaboration, employment, commission, sponsorship, endorsement, or institutional affiliation with any employer, client, school, partner organization, or other entity, and it does not represent any such party's views. No funding was received for this project except ordinary personal donations.

许可 / License

除非另有说明，本页原创文字、本站原创图片与本站原创图表采用 CC BY-NC-ND 4.0 发布。
转载时请保留原文标题、署名“JudgeDuck AI”、发布日期与原始链接；禁止商业转载、改写、摘编、翻译或基于原文创作演绎作品。

第三方商标、外部链接内容，以及文中另有标注的材料，不在上述许可范围内。

Unless otherwise noted, the original text, original images, and original figures on this page are licensed under the Creative Commons Attribution-NonCommercial-NoDerivatives 4.0 International License.