ds4.c 学习笔记

Appearance

学习日志 ​

Day 1 — 项目全貌与编译系统 ​

日期：2025-05-09 状态：已完成

今日摘要 ​

  • 理解了 Makefile 的完整构建链：make → 编译 .o → 链接可执行文件
  • 掌握了条件编译：-DDS4_NO_METAL 控制平台分支，__ARM_NEON 由编译器预定义
  • 读了 ds4.c 文件头：18 个系统头文件、模型架构常量 enum
  • 理解了 opaque type 封装：ds4_engine / ds4_session 前向声明
  • 学习了 DeepSeek V4 Flash 架构：43 层、284B 参数、256 专家 MoE（仅激活 6 个）
  • 关键洞察：专用引擎 > 通用框架——固定架构换来简化和优化

困惑与突破 ​

  • 困惑：为什么 make 里 ds4_metal.o 用 -fobjc-arc 而不是 -std=c99？一直以为 C 和 Objective-C 只是语法糖的区别。
  • 突破：.m 文件是 Objective-C，ARC 是 ObjC 的自动内存管理机制。C 和 ObjC 是两种不同的语言，编译器标志完全不同。ds4.c 用纯 C 是为了让引擎核心可移植，Metal 桥接单独用 ObjC 处理。

Day 2 — 内存管理与 mmap ​

日期：2025-05-09 状态：已完成

今日摘要 ​

  • 掌握了 xmalloc/xcalloc/xrealloc 包装器模式和为什么没有 xfree
  • 深入理解 mmap：MAP_SHARED（Metal 零拷贝）vs MAP_PRIVATE（CPU 避内核 bug）
  • 理解 xmalloc_zeroed 为什么不用 calloc（Darwin 零页 panic 风险）
  • 学习 posix_madvise 预取和 volatile 防优化技巧
  • 理解 __thread 线程局部存储和预分配 scratch buffer 模式
  • 追踪了完整加载链：open → fstat → mmap → cursor 解析 → weights_bind
  • 关键洞察：热路径零分配，所有内存操作集中到初始化阶段
  • 上游更新：model_open 新增 prefetch_cpu 参数，ds4_dump_text_tokenization() 传 false 避免遍历张量数据

困惑与突破 ​

  • 困惑：为什么 malloc + memset 比 calloc 更安全？calloc 不是天然初始化为零吗？
  • 突破：calloc 的"零页优化"让所有虚拟页指向同一个物理零页。推理时逐 token 写入 KV cache 会触发大量 page fault。在 macOS 上，大 mmap + 大量 page fault 叠加可能导致内核计数器溢出（kernel panic）。malloc + memset 在启动时集中触发所有 page fault，把风险控制在初始化阶段。

Day 3 — GGUF 二进制格式 ​

日期：2025-05-09 状态：已完成

今日摘要 ​

  • 学习了 GGUF 二进制文件格式：头部(24B) → 元数据表 → 张量目录 → 张量数据
  • 掌握了 cursor 游标模式解析二进制：has/read/skip 操作
  • 理解了 struct 对齐和 static_assert 编译时检查技巧
  • 学习了量化块格式：Q2_K(84B)、Q4_K(144B)、Q8_K(292B)、IQ2_XXS(66B)
  • 理解了 parse_metadata 的懒加载和 parse_tensors 的两步解析
  • 学习了 config_validate_model 的"早失败"验证
  • 关键洞察：懒加载 + 零拷贝字符串引用，避免不必要的内存分配

困惑与突破 ​

  • 困惑：cursor_string 返回的 ds4_str 没有拷贝数据，那 mmap 被释放后不就野指针了？
  • 突破：ds4.c 的设计是"加载后永不释放"——模型的生命周期覆盖整个进程。ds4_model 持有 mmap 映射直到进程退出，所以 ds4_str 指针始终有效。这是推理引擎常见的设计权衡：用内存常驻换取零拷贝。

Day 4 — BPE 分词器 ​

日期：2025-05-09 状态：已完成

今日摘要 ​

  • 学习了开放寻址哈希表：线性探测、2 的幂容量、& mask 代替取模
  • 掌握了 FNV-1a 哈希函数和动态数组翻倍增长模式
  • 深入理解 BPE 算法：字节编码→预分词→合并循环→查表得 token ID
  • 理解 GPT-2 字节级编码：不可打印字节映射到 256+ 码点
  • 学习了 Chat 模板组装：BOS + system + User + prompt + Assistant + 💭
  • 关键洞察：分词器是 LLM 的"眼睛"——文本必须先变成数字，模型才能处理

困惑与突破 ​

  • 困惑：BPE 合并顺序为什么是从最低 rank 开始而不是从最高？直觉上应该先合并最频繁的。
  • 突破：词表中的 rank 越小表示优先级越高。BPE 每轮扫描所有相邻对，找到 rank 最小（优先级最高）的对先合并。"最低 rank = 最高优先级"——这是训练时按频率排序编码的结果，rank 1 就是训练中出现最频繁的合并。

Day 5 — 采样策略 ​

日期：2025-05-09 状态：已完成

今日摘要 ​

  • 学习了数值稳定的 softmax（减最大值防溢出）
  • 理解 Temperature/Top-K/Top-P/Min-P 四种采样策略及组合
  • 掌握了插入排序维护 top-k 数组（O(n×k) 优于 qsort 的 O(n log n)）
  • 学习了 SplitMix64 PRNG 和均匀浮点转换
  • 追踪了完整采样链：session_sample → top_p_min_p → argmax/full_vocab
  • 关键洞察：采样策略在确定性和多样性之间找平衡，温度是核心控制旋钮

困惑与突破 ​

  • 困惑：为什么 softmax 要先减最大值再 exp？直接算 exp 再归一化不行吗？
  • 突破：浮点数溢出！比如 exp(1000) 是 infinity，整个计算就废了。减去最大值后，最大值变成 exp(0)=1，其他都是 exp(负数) < 1，不会溢出。这个"数值稳定性"技巧从 Day 5 一路用到 Day 8 的注意力计算。

Day 6 — 公共 API 与权重结构 ​

日期：2025-05-09 状态：已完成

今日摘要 ​

  • 深入理解 opaque type 封装：ds4_engine/ds4_session 前向声明
  • 掌握了结构体嵌套设计：engine → model + vocab + weights[43 layers]
  • 理解 weights_bind 按名称查找张量的绑定过程
  • 学习了 Transformer 每层的完整权重组成（注意力/FFN/HC）
  • 理解了超连接（Hyper-Connection）替代标准残差连接
  • 关键洞察：LoRA 低秩分解把 Query 参数减少 72%
  • 上游更新：ds4.h 新增 session rewrite API（ds4_session_rewrite_from_common），用于工具调用后检查点恢复；新增 ds4_dump_text_tokenization 独立分词调试
  • 上游更新（2025-05-15）：CUDA 新增 managed KV cache 支持——百万 token 级上下文时，KV cache 张量改用 cudaMallocManaged（统一内存按需分页），避免 DGX Spark 统一内存系统中 GPU 显存分配饿死 CPU 侧。普通上下文仍用 cudaMalloc 设备分配保持性能。通过 ds4_gpu_should_use_managed_kv_cache() 自动判断是否启用。

困惑与突破 ​

  • 困惑：LoRA 不是用于微调的吗？为什么 ds4.c 的 Q 投影用了 LoRA 结构？
  • 突破：LoRA 的低秩分解不仅用于微调——DeepSeek V4 把它作为模型架构的一部分。Q 投影从 4096 直出 32768 维太重了，拆成 4096→1024→32768 两步，参数减少 72%。这不是"微调补丁"而是"架构瘦身"。

Day 7 — 量化与矩阵运算 ​

日期：2025-05-09 状态：已完成

今日摘要 ​

  • 掌握了位操作提取 2-bit/4-bit 值的技巧
  • 理解了量化 block 布局：Q2_K(84B/256权)、Q4_K(144B)、IQ2_XXS(66B)
  • 学习了非对称量化策略：路由专家 2-bit、其他组件 4-8 bit
  • 理解了查表法在 IQ2_XXS 解码中的应用
  • 学习了 matvec 矩阵-向量乘法在推理中的核心地位
  • 关键洞察：284B 模型压缩到 81GB，在 128GB MacBook 上可运行

困惑与突破 ​

  • 困惑：2-bit 量化意味着每个权重只有 4 个可能值（00, 01, 10, 11），这怎么可能不严重损失精度？
  • 突破：关键在于"非对称量化"——不是所有权重都用 2-bit。路由专家用 IQ2_XXS（~2 bit）因为每个 token 只激活 6/256 个专家，量化误差被加权平均稀释。但 attention 投影和共享专家用 Q4_K 甚至 Q8_K 保持精度。这是"好钢用在刀刃上"的工程取舍。

Day 8 — 注意力机制 ​

日期：2025-05-09 状态：已完成

今日摘要 ​

  • 深入理解 Self-Attention：分数计算 → softmax → 加权求和
  • 掌握了 KV Cache 的必要性、滑动窗口和 memmove 实现
  • 学习了 Attention Sink：可学习 logit 参与 softmax 但不贡献 value
  • 理解了 MLA 压缩：KV cache 大小减少 64 倍
  • 掌握了混合注意力：raw (最近128) + compressed (历史) 并行
  • 关键洞察：压缩 KV cache 让百万 token 上下文在本地成为可能

困惑与突破 ​

  • 困惑：滑动窗口 SWA=128 意味着每个 token 只看最近 128 个 token，那 128 之前的信息不就全丢了？
  • 突破：两层机制保住信息：(1) 压缩 KV cache（MLA）把历史 token 压缩后保留，注意力时同时看 raw（最近 128）和 comp（历史压缩行）；(2) 43 层堆叠——每层看 128 个 token 的信息经过残差连接传到下一层，128×43 = 5504 的"感受野"。不是"传话游戏会传错"，而是"每层都在批注原始信息"。

Day 9 — MoE 混合专家 ​

日期：2025-05-09 状态：已完成

今日摘要 ​

  • 理解了 MoE 路由：哈希路由（前3层）和 top-k 路由（后续层）
  • 掌握了 √(softplus(x)) 路由激活和偏置选择/无偏权重的分离
  • 学习了 SwiGLU 激活：silu(gate) × up，以及 2-bit 量化时的 clamp
  • 理解了路由专家（2-bit）+ 共享专家（8-bit）的非对称精度设计
  • 计算了 MoE 效率：284B 参数中每次只激活约 0.1%
  • 关键洞察：MoE 用"宽度换深度"——参数多但计算少

困惑与突破 ​

  • 困惑：256 个专家只有 6 个被激活，剩下 250 个不是浪费吗？为什么不直接做 6 个大专家？
  • 突破：不同 token 激活不同专家——所有 256 个都会被用到，只是不是同时。43 层 × 6 专家 = 258 次专家调用，通过残差连接逐步积累。这好比一个医院有 256 个专科医生，每个病人只需要看 6 个，但不同病人看不同的 6 个。如果只有 6 个全科医生，每个人都要处理所有问题，反而更差。

Day 10 — RoPE 与推理循环 ​

日期：2025-05-09 状态：已完成

今日摘要 ​

  • 理解了 RoPE 旋转位置编码：通过 2D 旋转矩阵编码相对位置
  • 学习了 YaRN 扩展：混合外推和内插，支持 16× 上下文扩展
  • 掌握了 prefill vs decode：批量处理 vs 逐 token 串行
  • 理解了自回归生成的本质：每步输出成为下步输入，无法并行
  • 学习了 MTP 推测解码的原理和实现
  • 关键洞察：decode 的瓶颈是内存带宽，不是计算——每个 token 要读 81GB 权重

困惑与突破 ​

  • 困惑：RoPE 把向量做 2D 旋转编码位置，但为什么旋转后点积只依赖相对距离？这个性质是巧合还是设计？
  • 突破：这是 2D 旋转矩阵的数学性质：R(a)^T · R(b) = R(b-a)。两个旋转后的向量做点积，结果只包含角度差（相对位置），不包含绝对角度。RoPE 利用了这个性质，所以它不是"巧合"而是精心选择的编码方式。

Day 11 — 线程池 ​

日期：2025-05-09 状态：已完成

今日摘要 ​

  • 学习了 pthread API：mutex、condvar、create/join
  • 理解了线程池设计：worker 循环、generation 计数器、行范围分配
  • 掌握了 cond_wait 必须在 while 循环中的原因（虚假唤醒）
  • 理解了 __thread TLS 防止嵌套并行的机制
  • 关键洞察：主线程也参与计算，不浪费一个核

困惑与突破 ​

  • 困惑：为什么 pthread_cond_wait 必须放在 while 循环里？用 if 不就行了吗——signal 了就说明条件满足了？
  • 突破：POSIX 规范明确允许"虚假唤醒"——没有 signal 的情况下 cond_wait 也可能返回。而且多个 worker 被 broadcast 唤醒时，只有一个能拿到任务，其他必须继续等。while 循环保证每次醒来都重新检查条件，不信任任何唤醒。

Day 12 — HTTP 服务器 ​

日期：2025-05-09 状态：已完成

今日摘要 ​

  • 学习了 POSIX socket：socket/bind/listen/accept/recv/send
  • 理解了每连接一线程模型和单推理 worker 的架构
  • 学习了 poll() 处理慢客户端的 EAGAIN 场景
  • 掌握了 SSE 流式响应格式（data: + 双换行）
  • 关键洞察：推理串行但 HTTP 并发，用队列解耦
  • 上游更新（2025-05-15）：新增 /v1/responses 端点（OpenAI Responses API），支持 Codex CLI 等 Agent 工具。Server 日志新增 Responses API 标记和进度日志改进，方便调试多轮工具调用。

困惑与突破 ​

  • 困惑：为什么每个连接一个线程而不是用 epoll/kqueue 做事件驱动？C10K 问题怎么办？
  • 突破：ds4-server 的瓶颈是推理，不是连接。推理同时只能跑一个请求（GPU 独占），所以并发连接数在正常使用中不会超过几十个。每连接一线程简单可靠，是正确的工程选择。如果未来支持 batching，可能需要更复杂的事件模型。

Day 13 — KV Cache 持久化 ​

日期：2025-05-09 状态：已完成

今日摘要 ​

  • 学习了原子写入模式（tmp + rename 防崩溃）
  • 理解了 SHA1 缓存 key：基于 token IDs 而非文本
  • 掌握了磁盘 KV cache 文件格式（KVC 头 + 文本 + session payload）
  • 理解了四种保存触发和前缀匹配查找
  • 关键洞察：磁盘 KV cache 让长 prompt 的重复 prefill 从 10 秒降到 1 秒

困惑与突破 ​

  • 困惑：SHA1 基于 token IDs 而不是文本做缓存 key——为什么？文本不是更直观吗？
  • 突破：同一个文本可能被不同版本的分词器切成不同的 token。比如"你好"可能是一个 token 也可能是两个。用 token IDs 做缓存的 key 保证一致性——同样的 token 序列一定产生同样的 KV cache，不受文本编码变化影响。

Day 14 — API 兼容层 ​

日期：2025-05-09 状态：已完成

今日摘要 ​

  • 学习了手写递归下降 JSON 解析器：json_string/number/skip_value
  • 理解了 OpenAI vs Anthropic API 的格式差异
  • 掌握了 Tool Calling 流程：DSML ↔ OpenAI/Anthropic 双向转换
  • 理解了 Thinking 模式的三种级别和流式输出
  • 关键洞察：零依赖设计——JSON、HTTP、SHA1 全部手写
  • 上游更新：引入 rax 基数树实现工具调用文本记忆（精确回放模型采样的 DSML）；随机 tool call ID（/dev/urandom）；DSML 参数文本不再转义 <> &，只转义关闭标签；session rewrite 检查点恢复机制；JSON 解析器新增 JSON_MAX_NESTING 256 嵌套深度限制，防止深层嵌套 JSON 栈溢出 DoS
  • 上游更新（2025-05-15）：Anthropic live tool continuation——工具调用后保留 live KV 状态不变，下一请求通过 tool_use.id / tool_result.tool_use_id 匹配后只追加后缀 token，跳过完整 prefill。Responses API 同理：call_id 绑定 live KV 前沿，避免"工具调用后重建整个上下文"的停顿。核心思想：采样的 KV 状态是最高保真度的状态，客户端可见的协议对象只应"选择"这个状态，而不是强迫服务器重建它。详见 misc/RESPONSE_API.md 和 misc/ANTHROPIC_LIVE_CONTINUATION.md。

困惑与突破 ​

  • 困惑：手写 JSON 解析器为什么不用 state machine 而用递归下降？递归不是有栈溢出风险吗？
  • 突破：确实有风险——这就是为什么上游加了 JSON_MAX_NESTING 256 的深度限制。递归下降的好处是代码清晰、容易跳过不关心的字段（json_skip_value）。对于一个只需解析固定几个字段的推理服务器，递归下降的简单性远比完整 JSON 库更实用。

Day 15 — Metal GPU + 总结 ​

日期：2025-05-09 状态：已完成

今日摘要 ​

  • 学习了 ObjC 互操作：C 头文件隐藏 ObjC 类型
  • 理解了 Metal 计算管线：Device → Pipeline → CommandBuffer → dispatch
  • 学习了 Flash Attention 分块计算的思想
  • 理解了零拷贝 MTLBuffer：GPU 直接引用 mmap 内存
  • 用完整请求生命周期串联了 15 天的知识
  • 关键洞察：ds4.c 的哲学——垂直整合、专用化、零拷贝、热路径零分配
  • 上游更新：Metal shader 参数改为 char* 表达只读语义；encoder 可选 buffer 绑定零值占位满足 debug layer 校验
  • 上游更新（2025-05-15）：回退了 Metal Q4 expert tensor model views 的扩展（DS4_METAL_MODEL_MAX_TENSOR_BYTES 从 2 GiB 回到 ~672 MiB）。原因：该修复是由 AI 生成的且未经用户确认，上游坚持不接受未经人工验证的修复。

困惑与突破 ​

  • 困惑：Metal 可以零拷贝直接引用 mmap 内存，那 81GB 模型岂不是不需要任何显存？
  • 突破：Apple Silicon 的"统一内存"架构下确实如此——GPU 和 CPU 共享同一块物理内存。Metal 的 newBufferWithBytesNoCopy 只是创建一个指向 mmap 区域的 GPU 缓冲区描述符，不拷贝数据。但 GPU 执行时仍然需要把这些数据读入缓存层次，所以带宽是瓶颈。零拷贝省的是"额外的拷贝"，不是"不读数据"。

上游同步 — 2026-06-16 ​

子模块：vendor/ds4 从 c463029 前进到 e34a808（main 分支，43 文件、~27K 行新增）。本轮主线是 SSD streaming 的成熟化 + ROCm 后端并入 main。已把内容融入各章节：

SSD streaming（融入 Part 6 与 术语表） ​

  • 运行期路由热度淘汰取代纯 LRU：route_hotness[layer][expert] 表每 16 个 decode token 整表右移衰减一位，淘汰最低热度者。表是提示局部的（新 prompt 清零），但常驻缓存跨会话保温——区分"易变启发式"与"昂贵物化缓存"。
  • 三路 prefill 调度：极短走逐 token decode 式、中等走 selected-expert 批量、长 prompt 走全层 layer-major，按量化调阈值。
  • 缓存上限与 mlock 余量：显式 NGB 推理前封顶保持可锁定；mlock 失败时释放锁定余量、用实测可锁大小继续，而非硬塞可换页条目。
  • 混合精度路由专家（per-layer boosted quants）：检测提升层纳入 decode-span 兜底；slab 尺寸类启动时预置、对超大尺寸 freeze+reject 而非 last-writer-wins，防 slab 毒化。
  • streaming API 重构为传一张 ds4_gpu_stream_expert_table 表；CUDA/ROCm 不再是 stub，已实现 expert cache。
  • 新增 --ssd-streaming-preload-experts N（默认预热封顶 4096）。

ROCm / Strix Halo（融入 Part 6 新增 §4c） ​

  • ROCm 后端并入 main（此前只在独立 rocm 分支）。make strix-halo（别名 make rocm）用 hipcc 编 ds4_rocm.o + rocm/*.cuh（~18K 行），基于 rocWMMA。
  • 难点：补全 rocWMMA 内部头、扩大 GTT aperture（内核参数让 128GB 系统暴露 ~124GB GPU 可见内存）、避免 mixed IQ2/IQ4 GGUF 触发系统 OOM。详见上游 STRIXHALO.md。
  • 融合算子重构为可选后端钩子：ds4.c 不再写死 Metal/CUDA 分支，改为 ds4_gpu.h 上的可选钩子，新后端按需实现。

推测解码 / MTP（融入 Part 4） ​

  • 修复 --mtp-draft > 2 时验证器调错 topk_tensor 参数（单行取 top-N vs 每行取 top-1），导致接受了模型从未生成的 token。bug 只在深度 >2 暴露（深度 2 走专用 argmax 分支）。回归测试 --mtp-verify-depth 断言"每个提交 token 是其位置近似 argmax"这条不变量。

Agent / Serving（融入 Part 5） ​

  • 未关闭 <think> 内的工具调用恢复：检测到完整 stanza 开头时强制喂 </think> 前向恢复，而非重写已采样上下文。
  • edit 工具锚定 [upto] 匹配修复（old tail anchor 在 old head 之后找不到）+ 新增 ds4-agent edit 回归测试纳入 make test。
  • 终端 raw 模式状态在 quit/help/Ctrl+C 路径上正确恢复。

工程治理（融入 Part 1） ​

  • 新增 QA_BEFORE_RELEASES.md 发布前 14 节 checklist（覆盖 Metal/CUDA/ROCm/分布式/SSD streaming/KV cache/agent 等组合）。

困惑与突破 ​

  • 困惑：为什么运行期淘汰用"每 16 token 整表右移衰减"而不是逐 token 衰减或 LRU 时间戳？
  • 突破：逐 token 衰减每步要扫整张 [layer][expert] 表（43×256=11K 项），太贵；纯 LRU 又不适应工作负载切换。整表一次性 >>= 1 把衰减摊到 16 token 一个周期，开销恒定且小；新热点会因连续命中迅速盖过衰减变热，旧热点 16 token 后就降温——刚好匹配 decode 的局部性节奏。这是"精度换成本"的经典工程取舍：淘汰启发式不需要逐 token 精确，只需要相对热度排序大致正确。

上游同步 — 2026-06-28 ​

子模块：vendor/ds4 从 e34a808 前进到 80ebbc3（main 分支，10 commits、5 文件、~718 行）。本轮不是大特性，而是一波加固与稳定性修复：评测评分的假阴性、服务端 JSON 解析的边界、Metal 张量双重释放、ROCm 在 MTP/长 prompt 下的内存余量、分布式 + SSD streaming 的组合缺陷。已把内容融入各章节：

ds4-eval 答案提取器假阴性（融入 Part 5 §5） ​

四个会让"答对了却判错"的抽取 bug：

  • 行首代词/冠词遮蔽：find_answer_letter 取答案行第一个孤立有效字母，但 10 选项（A–J）题里行首的 I（I think...）、A（A careful...）、I'll 会被当成答案。改为跳过其后跟撇号或空白+小写字母的字母。
  • 被否定的干扰项抢答："It is not B, the answer is D" 被判成 B。新增 mc_letter_is_negated，回看前一个词是否 not/n't/rule out/eliminate/reject/except。线索集刻意做小、只回看答案标记后同行内——所以 "D, not B" 仍判 D（防住"取最后一个字母"的朴素错误修复）。
  • 整数答案行显示算式："m+n = 256+37 = 293" 被取成第一个加数 256。改为限定在答案行、且有 = 时取最后一个 = 之后的数。
  • 重评分级：regrade_trace_file 之前连 STOPPED/SKIPPED 的残缺输出都评，虚增计数、制造假漂移。抽出 regrade_case_outcome，只评 PASSED/FAILED，其余计 not_graded。
  • 全部由 --self-test-extractors 无模型用例锁定。

服务端 JSON 解析加固（融入 Part 5 §3，issue #405） ​

OpenCode 这类工具密集型请求压出来的三处边界：buf_reserve 翻倍循环在 SIZE_MAX 附近回绕（防回绕检查）；畸形代理对（\uD83D 无合法低代理）吃掉相邻转义（回退指针归一化为 U+FFFD）；新增 32 工具 × 24 轮的 tool-heavy 压测 + 代理对单元测试。顺带把 stop-list 的 strstr 返回值改 const，保 -Wall -Wextra 无警告。

Metal 张量双重释放守卫（融入 Part 6 §Metal 与 术语表，issue #404） ​

ds4_gpu_tensor 句柄是 retained ObjC 对象，重复释放会让 __bridge_transfer 释放已释放对象 → malloc corruption。新增一张活跃句柄集合（开放寻址 + 墓碑 + 70% rehash），由 g_tensor_mu 保护，free 前先查集合、不在则拒绝。诊断计数器也走同一把锁。这是个少见的"真用上墓碑"的开放寻址实战用例。

ROCm 稳定性（融入 Part 6 §4c） ​

  • 更安全的 prefill 默认（#387）：移除 ROCm 对长 Flash prompt 的 8192-token 工作区默认，改用普通 4096——在 Strix Halo 上 IQ2 全缓存 + 8192 会耗尽 HSA 分配空间。
  • MTP 启动驻留（#425）：第二个 GGUF 映射（MTP）注册时，range 查找改为先用已备 range 表和 fd 缓存、最后才回退无界拷贝，让主模型与 MTP 共存；并在多模型时禁用可选的 Q8→F16 扩展缓存，给 session/context 张量留余量。

分布式 SSD streaming 层切片（融入 术语表、术语表、Part 6 §4a） ​

ds4_session_eval_layer_slice 在 streaming + 分布式切片组合下，流式映射设一次管不了整层栈——映射是 per-layer、按 decode/prefill 模式分的。改为每编码一层前重新映射、每层包自己的命令边界，让两者能正确叠加。

困惑与突破 ​

  • 困惑：为什么行首代词遮蔽只在 10 选项题里出问题？4 选项（A–D）的题里 I think it is D 不也满足"行首孤立有效字母"吗？
  • 突破：满足，但 4 选项的有效范围只到 D，I 根本不在有效字母范围内，find_answer_letter 的 c >= 'A' && c <= max_answer 直接跳过它——所以无害。只有当干扰字母本身落在 [A, max_answer] 区间内（10 选项 A–J 时 I 进入了范围），它才有机会"假装是答案"。这是"边界条件 + 数据集规模"共同决定的可触达性：bug 一直存在，但只有 24 道 10 选项用例能触发它。

上游同步 — 2026-07-29 ​

子模块：vendor/ds4 从 80ebbc3 前进到 54b36ed（main 分支，51 commits）。本轮是项目自开源以来最大的一次扩张——DwarfStar 从"DeepSeek V4 Flash 专用"走向多模型 + 多机多卡：新增 GLM 5.2 整条架构支持、张量并行（Mac RDMA + CUDA 两种形态）、会话批处理、DSpark 推测解码、多 GPU 放置，外加分布式 CLI/文档统一与一批正确性修复。已把内容融入各章节：

GLM 5.2 —— 第二种模型架构（融入 Part 3 §5、Part 6 §4l） ​

引擎加载时按模型族分发（ds4_engine_is_glm_dsa()）。GLM 同为 dense + 路由 MoE，但量化布局不同：dense/控制张量走 Q8/F32，路由专家 gate/up 支持 Q2_K/Q4_K/Q5_K、down 支持 Q2_K/Q4_K/Q5_K/Q6_K。model_id（0=Flash）做缓存兼容性稳定 id。MTP 块在主 GGUF 内（--glm-mtp），不支持方向性引导/--power<100/外部 --mtp 文件。ROCm 侧为 GLM 写了 ~4.2K 行调优（rocm/ds4_rocm_glm.cuh）。GLM 的 GGUF 由 download_model.sh glm-* 目标下载。

张量并行（融入 Part 6 §4j、术语表） ​

新增 ds4_tp.c/h（~2.2K 行）。与流水线并行根本不同：两机在同一 token 上同时切层内工作量，在每层 ATTN/FFN 同步门交换 16–24 KB f32 部分和（不量化）。形态一是 Mac-to-Mac RDMA over Thunderbolt（Leader 镜像每次 sync/eval，lockstep）；形态二是单机 CUDA --cuda-tensor-parallel（偶数卡，先列 home 再列 partner）。握手交换 ds4_tp_identity，门调度两边必须一致（DeepSeek 每层 ATTN+FFN，GLM 只在稀疏层发 FFN 门）。不能与 SSD streaming/分布式/MTP/CPU 并存。

会话批处理（融入 Part 5 §6、术语表） ​

--batched-session N 预分配 N 个常驻 KV session，就绪 decode 步打包进同一内核。精确性保证：无原生批处理内核时按固定顺序跑受影响行，返回与逐 session 单跑一致的 logits。per-backend 行为：Metal Flash 原生共享专家+QKV 批处理、GLM 有序兜底、CUDA TP 原生 decode+混合 prefill/decode。批处理激活时 MTP 禁用。顺带修了 session 恢复与并发 decode/prefill 的竞态（恢复完成前不参与调度）。

DSpark 推测解码（融入 Part 4 §5、术语表） ​

DeepSeek 为 Flash 发布的辅助草稿模型，读主模型隐藏状态、一次提最多 5 个 token，主模型验证后提交接受前缀。取代旧单阶段 MTP（不叠加），用 --mtp 传 ~5.6 GiB 支持文件、--dspark 选运行时。Flash 专用（PRO 暂不支持），默认置信度门 0.9，采样解码不用其提议。不加速 prefill，可预测续写（代码）收益最大。

多 GPU 放置与层打包（融入 Part 6 §4k） ​

三个新模块：ds4_gpu_args.c/h（--gpu-vram/--gpu-devices 统一解析，auto/0/显式预算）、ds4_layer_pack.c/h（层序列单调连续放进各卡预算，溢出到 CPU 层级，无"层太大"错误路径）、ds4_gpu_mgpu.h（ds4_gpu_tensor 带 device_id、g_gpu[16]、跨设备 P2P 拷贝、TP 归约 ds4_gpu_add_xdev_tensor）。引擎入口 ds4_engine_create_with_gpu_config() 接受可选 ds4_gpu_config。

分布式统一 + 工程治理 ​

分布式推理的 CLI 和文档统一进 README（流水线并行 PRO Q4 双 Mac Studio、网络链路对比、分布式协议概述）；README 整体重写（DwarfStar logo、多模型定位、AI 协作披露）；make cuda-spark/cuda-generic/strix-halo 目标。正确性修复：ROCm IQ2 路由专家 LUT 初始化与 MoE 正确性、批处理 server session 恢复竞态、PRO streaming/采样正确性、release 构建零警告。

困惑与突破 ​

  • 困惑：张量并行为什么"路由专家按连续一半切分、dense/注意力/共享专家却复制"？既然要 50/50 省内存，为什么不把所有东西都切一半？
  • 突破：因为路由专家占模型绝大部分空间，却高度稀疏——每个 token 只激活 6 个（见 Day 9 MoE）。把它们对半切，容量直接翻倍，且每个 token 的 6 个专家要么在本机要么在对方——路由内核按 ownership 只读自己那一半，永不跨机读专家。而 dense/注意力/共享专家体积小、且 lockstep 图里两边每层都要用，复制它们的成本远低于"每次同步门都协调切分"。这和 SSD streaming（路由专家大部分时间在睡觉、可缓存淘汰）是同一条洞察的两种应用：MoE 的稀疏性是 ds4 一切"容量魔法"的物理基础。

上游同步 — 2026-08-08 ​

子模块：vendor/ds4 从 54b36ed 前进到 b030961（main 分支，86 commits）。本轮是 Flash 0731（DS4F）落地 + CUDA 性能大跃进：新增 MXFP4 量化格式（CUDA/Metal 原生推理 + 无损 GGUF 转换）、DeepSeek V4 Flash 0731 检查点成为 L40S 默认、把 batched-serving fork 的 vendored MMQ 与 decode-island CUDA graph 接入上游，外加 DSpark 调优、Metal mlock 回归修复、ROCm 稳定性、Blackwell sm_120a 支持与一波 agent/JSON 加固。已把内容融入各章节：

MXFP4 —— 微缩浮点量化（融入 术语表、Part 6 §4m） ​

新增 OCP Microscaling FP4 格式 block_mxfp4（17 字节 / 32 权重 = 4.25 bit）：1 字节 E8M0 共享指数 e（幂次缩放 2^(e−127)，e=0 视为零）+ 16 字节 nibble（每个 4-bit 索引一张 E2M1 码表 {0,±0.5,±1,±1.5,±2,±3,±4,±6}）。反量化 w = e8m0(e) × codebook[nibble]，GGUF 张量类型 39。关键区别于 Q4_K：MXFP4 是真浮点（硬件原生 FP4）+ 每 32 元素共享指数，可直接映射 GPU 张量核的 block-scaled MXFP4 指令，不必先反量化。725b084 做无损转换——DeepSeek 原生 MXFP4 检查点按字节进 GGUF，无 dequant→requant 损耗。这正是 Flash 0731 的原生格式。

DeepSeek V4 Flash 0731（DS4F）+ 基准刷新 ​

新检查点 0731（代码里称 DS4F）是 MXFP4-native Flash。7b55707 把 L40S server 默认指向 Flash 0731 MXFP4；ae504c9 更新下载脚本与 DSpark 支持文件；fixture 按检查点版本化（b7e9f00）。基准表用 ds4-bench（Promessi sposi 输入、2048-token 阶梯、128 贪心 token）重测，完整曲线在 speed-bench/m5_max.csv / gb10.csv：

机器后端上下文PrefillGeneration
M5 Max 128GBMetal2048790 t/s39.4 t/s
M5 Max 128GBMetal65536399 t/s27.6 t/s
DGX Spark GB10CUDA2048826 t/s18.1 t/s
DGX Spark GB10CUDA65536823 t/s13.8 t/s

DS4F 还暴露了一个行为差异：新检查点对锚定 [upto] 编辑吃力（b030961），故 agent 的 edit 默认改回精确 old/new 替换，[upto] 提示与自动插入改为 --edit-upto 显式 opt-in。

CUDA 性能大跃进（融入 Part 6 §4m） ​

三个大件把 batched-serving fork 的工作接入上游：

  • Vendored MMQ prefill tier（0f62c62）：从 Entrpi/ds4 fork 引入 cuda/mmq/（llama.cpp 的 Multi-Marlin Quantization 内核），只接 RAW-layout 层——n_tok≥2 的 dense Q8_0 GEMM + IQ2_XXS gate/up + Q2_K down 路由 MoE 管线；decode（n_tok==1）不动。代价：FP32 归约顺序变 → prefill logits 在 ULP 尺度漂移，故仅在 quality 关闭时启用，DS4_CUDA_MMQ=0 回退。配套 6f087f4：make cuda-spark 之前空的 CUDA_ARCH 让 nvcc 退到默认架构，把 MMQ 张量核路径整个编掉（实测 39 vs 423 prefill tok/s！）。
  • Decode-island CUDA graph 捕获（e50104f）：把每个 decode 层的两个位置无关"岛屿"（hc-pre 前缀 + 注意力输出到层尾的尾部）按激活缓冲变体捕获一次、后续 token 直接回放。捕获/回放复用现有 phase 机制，无重复编码路径；首次见 key 走 eager 预热、二次捕获、之后回放；失败则退役重编码，不丢 token。TP/多级/SSD streaming/debug/hash-router 层都自动 disqualified 保持 eager（DS4_CUDA_DECODE_GRAPHS=0 关闭）。DGX Spark GB10 上 decode +1-2%——说明 2k 上下文下 launch 开销不是串行 decode 的主成本，捕获层是为后续覆盖位置相关中段做的地基。
  • Blackwell sm_120a（b1f8893）：RTX 50 / PRO Blackwell 编为 compute_120a/sm_120a，这样才接受 block-scaled MXFP4 指令；兼容 sm_120 与 sm_120a 两种写法。
  • 一批 HMMA / indexed-attention 内核（token-tiled 注意力 cfaee47/bff6f5e/1728313、indexer scorer 去 bank-conflict c784026、SM121 MXFP4 indexer 56ec892、HC RMS norm 折进 FP16 投影输入 574ccbb 等），小 Q4 TP 批次加速（1464c0f）、批处理 server prefill quantum 可配（8579ec8）、GLM 批处理内存计入放置（6537ba0）。
  • CUDA TP server 托管重启（1c5e24b）：run-nvidia-tp-server.sh 让 --cuda-tensor-parallel server 崩溃后自动重启；顺带修了启动与 toolkit 检测（108eb3a），并在评估前拒绝不支持的 GLM TP 专家布局（6747e77）。

Metal：mlock 回归修复 + Metal 4 prefill（融入 Part 6 §2/§4a） ​

GLM 5.2 重构时把 streamed-expert 的 mlock pinning 全删了——mlock 调用、slab lock、预算上限全归零。SSD streaming（模型 >> RAM）下 OS 在内存压力下驱逐专家缓存，每个 token 都从磁盘重读，Flash decode 回归约 2-3× 变慢、首 token ~6s（#532）。3ce6777 把 pinning 适配重构后的 slab 分配器移植回来：整缓冲 mlock、per-slot fill 锁 / relief 解锁、routed-memory 预算上限优雅降级、失败时 cap-on-failure、压力下解钉最冷的 ~10%。M5 Max 128GB 实测从 0.3-2.5 → ~7.0 gen t/s，反超之前基线。另有一批 Metal 4 prefill 加速：路由 MoE prefill（532ec8b）、indexed prefill（222b2cb）、Q4 attention 输出 prefill（96c3ba4）、partial Q4 prefill tiles（d69a017），以及 streamed prefill maps 与专家内核一致性（d14ce35）、从映射 prefill 层播种专家缓存（4893e0c）。

DSpark 调优（融入 Part 4 §5、术语表） ​

  • 确定性调度默认开（4591cb1）。
  • 默认置信度门从 0.9 降到 0.7（a51e6ec）——放更多后缀过验证，换更高接受率。
  • 贪心同一性修复（af80694 / 7fb2830）：partial-accept 捷径之前保留了验证器批的 compressor 前沿，可能改变后续贪心 token；现在恢复快照、把 partial accept 像 full accept 一样走普通单 token 路径回放，保证推测验证不留数值不同的 compressor 状态。顺带：final hidden 工作留在 stage command buffer（f284165）、严格 Flash 捕获跳过无 compressor 的前两层（422ceb1）。

ROCm 稳定性（融入 Part 6 §4c） ​

修复 DeepSeek 在 ROCm 上的一串稳定性问题：IQ2 SSD prefill 死锁（2f49b27）、Q4 SSD 专家 staging（9e1988b）、大模型 arena 收紧（6882a0f）、prequant decode 内核恢复（1e8f16c）、两 token MTP 验证优化（d250a7c）。

Agent / JSON 加固（融入 Part 5 §3） ​

一波解析器与协议加固（多数有 issue 编号）：

  • JSON 双重释放（#539，8e49a5c）：解析前清空字符串输出，畸形重复字段不能再留下已释放指针。
  • 先解析成功再替换（#433，5c318ed）：重复字段先进临时存储、解析成功后才释放旧值，覆盖 OpenAI/Anthropic/Responses/tool schema。
  • NaN 拒转整数（#541，8340f35）：请求字段转 int 前先夹住非数与负值，避免未定义浮点转换。
  • 不可能的模型文件尺寸（#542/543/545，c7689db）：safetensors/GGUF 长度字段按文件剩余字节封顶、交叉校验 FP8/FP4 张量存储、分配前拒绝不可能的元数据与张量计数；配套 size 检查溢出安全（a968c08）、超大快照分配前拒绝（#546，f02d79c）、agent cache 字符串按剩余文件封顶（#544，0fa15c6）。
  • 工具历史线性校验（#547，a169cff）：边正扫边建"最近一次调用"查找，大 Responses/Anthropic 历史不再触发反复全历史搜索。
  • 未终止 reasoning 路由（#509/665，9735ad8）：截断的 DeepSeek/GLM reasoning 在非流式响应里导进 reasoning_content，不污染正文。
  • 重复 reasoning 不进流式答案（#678，fe2d3b0）、未关闭 reasoning 内恢复完整工具调用（#675，51a1c14）：等完整调用、保留前面散文为 reasoning、结构化返回、协议文本不进流式输出。
  • 重复 GLM prompt 复用（#668，7694112）：incoming GLM prompt 恰是 live session 开头时，回退到边界前一 token 重评估，不重建长 prompt。
  • 回归测试（355da75）+ critical-fix 发布门（7adb2a4）、速度回归门（0dd0d36）锁定上述修复。

KV cache（融入 术语表、Part 5 §2） ​

24fa85e（#673）：磁盘 KV 检查点成功加载后不再被 cold-prompt 保存上限误删——更深的 cold / continued 检查点首次恢复后保留，磁盘预算驱逐仍负责回收文件。

困惑与突破 ​

  • 困惑：DSpark 把默认置信度门从 0.9 降到 0.7，接受的后缀更多，但单次验证的算力是固定的——这不是在"赌"更多 token 通过吗？降低门槛难道不会让被拒后缀的算力白花？
  • 突破：DSpark 的验证不是逐个 token 决定接受/拒绝，而是对一个 5-token 块取被接受的前缀。门限切的是"后缀从哪里开始不可信"——0.9 太严，把本可接受的后缀整段丢掉、退化成只接受 1 个 token，验证算力没被充分换回 token；0.7 放宽让前缀更长，单次验证换回的 token 期望值更高。代价是被拒后缀的草稿算力确实"白花"，但草稿模型小、读主模型隐藏状态，单次提议远比一次主模型目标解码便宜——只要平均接受长度上升的收益超过偶尔多拒的草稿成本，门槛就该下调。这是"验证算力固定 → 最大化每次验证换回的 token"的直接优化，配合贪心同一性修复（保证宽松接受不污染 compressor 状态）才安全。

上游同步 — 2026-08-14 ​

子模块：vendor/ds4 从 b030961 前进到 84cc882（main 分支，30 commits、33 文件、~11.9K 行新增）。本轮主线是 Metal M5 Max 精确解码融合战役：把 decode 从 39.4 推到 ~45 t/s 的同时保持逐位精确，并把融合移植回 M1-M4；外加 DSpark 语义升级（直接提交验证器状态 + 登陆 ROCm）、服务器的客户端断开取消，以及长上下文正确性修复。已把内容融入各章节：

Metal 精确解码融合战役（融入 Part 6 §4n、code-walkthrough 追踪 9） ​

  • 方法论："exact" = 与旧路径的全词表 logits 逐位相同 + token 选择一致，由两个新 A/B 基准（make metal-decode-schedule-bench / metal-prefill-variant-bench）强制：交替变体顺序与 session 分配、毒化 logits 缓冲，不逐位相同即 abort。
  • 动机：decode 已 ~97.5-98% GPU-busy，瓶颈是 dispatch 次数（每层 ~25 次 × ~3.9µs）——融合相邻内核，一次 dispatch 干掉多次启动开销。
  • 成果：十余个融合（HC norm+mix、四路 F16 压缩投影、q_a/kv pair+quad、emit 路径 finalize、压缩器精确 softmax+product、聚簇 HC 生产者、打包 split-K FA 归约（+3.95%，消除每层 4.2MB device partial 往返）、并发全 FFN 等）。M5 Max q2 decode 39.39 → 45.25 中位 / 45.47 稳态（+14.9%）；移植回 pre-M5（M3 Ultra）41.08 → 44.18（+7.55%）。
  • 门控自动化（0ad494e）：删 ~76 个 DS4_METAL_ENABLE_* opt-in 与 ~27 个废弃 M3 别名，融合按设备（ds4_gpu_device_is_m5/pre_m5_apple_silicon()）+ 形状默认启用；DS4_METAL_DISABLE_* 只作回滚控制。5186e2c 给非 Apple 构建补 stub 保编译。
  • MXFP4 exact 加速（Ivan Fioravanti，fd2d5e0）：half 域缩放、KV RoPE 折进 FP8 存储、逆 RoPE 折进 FA reduce，M3 Ultra 15 轮 prefill +10.9% / decode +9.5%；MXFP4→half 查找表（metal/generate_mxfp4_half_lut.py 生成 256×16 精确 binary16 表，make check-mxfp4-half-lut 钉 SHA-256 防漂移，decode 步 +4%）融入 量化术语表。
  • 正确性：0658dac 修两个长上下文 bug——逆 RoPE 融合的全局 consumed 标志被后续 indexed-attention 层误读（改局部 per-layer 跟踪）、kernel_dsv4_kv_rope_fp8_store_f32 的 threadgroup 屏障对设备内存写入不足（改 mem_device_and_threadgroup）；023614e 删除输出顺序不可复现的 kernel_topk_stream512，永远走确定性 argsort+merge。
  • 负面结果入档：Q8_MV_ROWS=4 的 43.7 t/s 实为跳行损坏；M5 上 NAX n_tok=1 matmul 不逐位精确被拒；只信进程内交替对比（跨时间窗机器漂移 ~10%）。

DSpark：直接提交验证器状态 + 三后端（融入 术语表、Part 4 §5） ​

  • 0e89a0e：接受后不再"回滚快照 + 逐 token 重放"，直接采用批量验证器的状态（full accept 直提；partial accept 用前缀快照槽位直提，失败回退重放；TP partial 仍重放）。与逐 token 解码跑同一张图、浮点分组顺序不同——字节同一性不再保证（非降精度），--dspark-strict 保留逐字节复现模式；QA 从"byte-identical"改为记录首个差异、直接提交数、回退数。
  • 42033ee：Metal 上 DFlash 验证小批量每 4 层 flush 一次（host 编码与 GPU 执行重叠）；置信度默认改为 Metal 0.6 / CUDA+ROCm 0.7。
  • 84cc882：补齐 ROCm 草稿/验证内核（非因果批量注意力、markov argmax、attention launch），DSpark 覆盖 Metal/CUDA/ROCm 三后端。

服务器：断开取消 + 工具调用恢复（融入 Part 5 §3b/§3） ​

  • e9ded97（+测试 8a703b6）：客户端断开后不再空算——client 线程每 100ms 探测连接（零超时 poll + 排干 pipeline 字节看 EOF），断开即经 ds4_session_set_cancel() 取消在途生成（与 agent Ctrl+C 同一套合作式中断机制）。
  • 0ead8a8：未关闭 thinking 内的截断工具调用也走恢复——检测改三态（NONE/STARTED/COMPLETE），STARTED 在生成结束时交给 try_repair_dsml 补关闭标签；完整 stanza 照旧前向恢复。
  • 3196149：工具 schema 按规范词法重渲染（键插入序 + Python 分隔符 + 解码 UTF-8），语义相同的请求得到字节相同 prompt 前缀，保 KV cache 命中。

基准 ​

e0c63d9 把 M5 Max / GB10 的 q2 全曲线表固化进 README（M5 Max 2k：790 prefill / 39.4 gen；GB10 2k：826 / 18.1，65k 仍 823 / 13.8）——与 §4n 战役数字（45.25）是两份不同契约，笔记里已分别标注口径。

困惑与突破 ​

  • 困惑：既然要求逐位精确，浮点运算不是有结合律问题吗？两个内核做同样的加法，顺序不同结果就可能差一个 ULP——"bit-exact 融合"怎么做到的？
  • 突破：融合内核不是"重新实现同样的数学"，而是把旧 dispatch 的归约树拓扑原样搬进新内核——相同的线程数、相同的分组、相同的舍入点，只是把"多次 dispatch 之间的显式依赖"变成"一次 dispatch 内的 threadgroup 协作"。比如聚簇 HC 生产者明确"保留精确 1024 虚拟线程 RMS 树"，打包 FA 归约"保留原 32-lane 归约树"。这就是 A/B 基准敢要求逐位相同的底气：算术没变，变的只是调度。反过来，一旦某条路径真的改变了运算顺序（M5 上的 NAX n_tok=1 matmul），哪怕更快也被拒绝——"exact" 是合同条款，不是形容词。

上游同步 — 2026-09-01 ​

子模块：vendor/ds4 从 84cc882 前进到 b0982a1（main 分支，~100 commits、~31K 行源码新增）。本轮三条主线：GLM 5.3 整条架构（KDA 递归线性注意力 + 视觉 + 全量检查点 SSD streaming）、DeepSeek V4 Flash 视觉（Vision-Exp）、ROCm 大扩张（MXFP4 专家内核 PR #656 + gfx1151 调优战役把 prefill 从 ~226 推到 290.79 tok/s）；外加 DSpark 采样语义反转与服务器多模态 KV 安全。已把内容融入各章节：

GLM 5.3 —— 混合递归/稀疏架构（融入 Part 3 §6、KDA 术语表、Part 6 §4l） ​

46 层 / 4096 维 / 154880 词表 / 288 专家取 8+1 共享 / 4 路超连接 / MTP 内建。架构核心是层混合：ds4_glm53_layer_is_kda(il) 判 il % 4 != 3——3/4 的层是 KDA（Kimi Delta Attention，门控 delta 规则递归线性注意力，没有 KV cache，历史卷进每头一个固定大小状态矩阵），每第 4 层是完整 DSA 稀疏注意力 + pool-4 索引器负责精确检索。decode 数学五步：深度 4 短卷积 → 逐通道衰减 d = exp(lower_bound·sigmoid(exp(a_log)·(gate+dt_bias))) → L2 归一化 → delta 规则 H = H·d; δv = (v − H·k)·sigmoid(β); H += k ⊗ δv; o = H·q → RMSNorm + sigmoid 输出门。配套：Q2/Q4 工件把 KDA dense 全升 Q8_0（递归路径不容量化噪声累积）、glm53-full-q2 全量检查点 streaming（常驻层区间整层映射 + glm_graph_memory_guard_budget 自适应预算）、零 RoPE 尾巴的 FP32 紧凑 KV（DS4_GPU_GLM_COMPACT_CACHE_F16）、M3 Ultra BF16 融合 QKV（ds4_gpu_glm53_matmul_bf16_qkv 一内核产 q/k/v，decode 24.74 t/s）、续 prefill 切片修复（639d4eb）。下载：glm53-q2（~90 GiB）/ glm53-q4（~178 GiB）/ glm53-fp8（~305 GiB，仅打包）/ glm53-full-q2。

视觉管线 —— 两族多模态（融入 Part 3 §7、视觉术语表、Part 5 §7、Part 6 §4o） ​

新模块 ds4_image.c + vendored third_party/iris（头文件-only JPEG/PNG 解码，EXIF 方向，SHA-256 指纹）；编码器 GPU-only，权重是 --vision 侧车 GGUF（Apple 直接映射 / CUDA/ROCm ds4_gpu_set_aux_model_map_range）。GLM 5.3：Qwen-VL smart resize（1568 像素/token、16..8000 预算）、24 层 ViT + QK-RMSNorm + 2×2 merge 投影，进上下文靠"占位 token + 浮点嵌入覆盖"（glm_vision_overlay_prepare）。DeepSeek Vision-Exp：硬 384 token 预算（压缩比 3 网格）、32 层 ViT（中间结果舍入 BF16 对齐参考数值）、合成哨兵 token + 行对交织，且语言模型带 visual_router_bias[43]——视觉 token 走不同的 MoE 专家分布。多模态 KV 安全是本轮的服务器侧重点：ds4_session_vision_state_matches（指纹+跨度全等才复用）、恢复起点推进到 live 前沿、磁盘 KV 拒收视觉状态（98b9135/b0982a1）。API 契约：三协议内联 base64 图（png/jpeg）、≤16 张、64 MiB body、路径与 URL 拒绝。

DSpark 采样语义反转（融入 Part 4 §5、DSpark 术语表） ​

769a8ba 实现分布精确的随机推测解码（Leviathan/Chen：点质量草稿 q、sample_build_probabilities() 物化用户过滤链后的 p、接受 = min(1, p)、拒绝按残差重采样——输出分布与直接采样目标严格一致）；8d683d6 随即反转默认为机会式温度采样：边界 token 按请求温度正常采样，DFlash 贪心后缀与目标贪心路径一致即提交（可预测代码 temp 1 下 M5 Max +8%）；精确模式降级为 --mtp-exact-sampling（未设时置信度门 0.8、草稿 ≥2）。正确性由 10 万次 Monte-Carlo 频率断言锁定（误差 < 0.006）。另有 ROCm 的 seed-plus-five 验证器融合（DS4_DSPARK_SEED_BATCH，贪心一致才启用）与零门限探测跳过。

ROCm：MXFP4 内核 + gfx1151 战役（融入 Part 6 §4c、量化术语表） ​

  • MXFP4 路由专家内核（PR #656）：GPU 上解码 + Q8_K 激活 + dp4a；V_PERM_B32 一指令展开 4 码字、lane 对分块合并访存、lo/hi 整数累加器（重结合但逐位一致）、gate+up 融合装载；固定 8 次全展开让 gfx1151 代码生成 99→83 VGPR / 12→16 waves；(tile, row-block) L2 友好网格序让专家权重从 DRAM ~只流一次（−10.3%，down 侧同交换负收益即回退）；占用率变体全部逐位精确 + env 门。make test-mxfp4-rocm 按 Flash 真实形状测。
  • gfx1151 调优战役（~226 → 290.79 tok/s prefill）：kernel lab 工作流——cuda/mmq/test/proto_*.cu 独立原型基准（cuBLASLt best-of 做上限参照），赢家经 ds4_rocm_gfx1151_flag() 默认接入，数字逐轮记入 speed-bench/gfx1151-prefill-results.md。主题：80-key WMMA 注意力 tile、IQ2 符号掩码缓存、成对 gate/up tile、small-M F16 GEMM 模板化、tiny-Q8 verifier（24 VGPR 胜 72 VGPR 占用率悬崖）、rocBLAS 解方案钉扎 + 被拒后永久回退。vendored cuda/mmq/vendors/hip.h 让 CUDA MMQ 在 HIP 下编译（DGX Spark GLM prefill 也因此受益）。
  • GLM 5.3 on ROCm：rocm/ds4_rocm_glm.cuh 增至 ~5.3K 行，KDA 三步 prefill + decode 内核全套；fea665c 修块归约缺 __syncthreads() 的共享缓冲竞态。make test-rocm 让 ROCm-only 主机跑核心回归。

服务器 / 工具（融入 Part 5 §2c/§7/§4） ​

ignore_eos（要求 temperature 0，ds4_session_argmax_ignoring_eos 跳过停止 token——评测用固定长度贪心）；live steering 库级 API + REPL /steer（不重建 KV 调强度，87495f6；GLM 5.3 引导 8adb221）；agent 接受 --rocm 后端、--vision 启用 view_image 工具；流式终端折行修复（DECSC/DECRC 不保留自动换行 pending 标志，04478e0）；模型动物园加 V4 Pro 0813 quant（pro-q2-imatrix ~430 GB）与 ds4f-vision-*（含专用 DSpark，不能混用 0731 版）；PRO 量化工具修复（compress_ratios 改从 HF config.json 读、Q2_K/Q4_K 精化循环终止条件）。

困惑与突破 ​

  • 困惑：GLM 5.3 敢让 3/4 的注意力层完全没有 KV cache——把历史卷进一个固定大小的状态矩阵，长上下文的精确回忆怎么办？不怕"传话游戏传丢"吗？
  • 突破：这正是"混合"二字的用意。KDA 的 delta 规则不是朴素累加：H += k ⊗ (v − H·k) 先探测状态里已记住的内容、只写残差回去，比普通线性注意力更能改写旧关联——但它终究是有损压缩，固定状态装不下任意长的精确历史。所以每第 4 层保留完整 DSA 稀疏注意力 + pool-4 索引器：递归层负责 O(1) 内存的"大意"，全注意力层周期性地做"无损校验点"。这和 Day 8 学的混合注意力（raw 窗口 + 压缩历史）是同一个设计模式的两档：MLA 把历史压缩了但仍按 token 存行，KDA 干脆把历史归零进状态、再用稀疏的全注意力层兜底。省下来的内存是实打实的：46 层里 35 层的 cache 行从"随上下文线性增长"变成"常数"。

上游同步 — 2026-09-08 ​

子模块：vendor/ds4 从 b0982a1 前进到 f62ca29（main 分支，96 commits、~22.6K 行新增）。本轮一条主线压倒其余：两 Mac Metal 张量并行的延迟工程战役（macm5us/mac-tp-1-sep-2026 合入，decode 16.8 → 47-48 t/s、prefill ~94 → ~850 t/s @2048）；配套 Metal prefill 战役、GLM SSD streaming 成熟化、DSpark 种子批处理 + 六 token 验证、ROCm gfx1151 第二轮（PR #992）、agent/服务器工具调用加固与对话前缀条件化（--prefix-file）、ds4-eval hard 套件。README 拆分成 12 篇 docs/ 专题指南。已把内容融入各章节：

双 Mac TP 延迟工程（融入 Part 6 §4j、张量并行术语表） ​

协议升 version 10（新帧 RDMA_WARM/POSTED/GLM_MTP；GLM TP 推测周期有显式命令与共享 token 上限——两端必须一起升级）。一周 ~25 个提交全是微秒级抖动打磨：轮询门（CPU→GPU 事件唤醒 27-43 µs，改 16 区 flag 环轮询）+ 校验和 flag（flag 行比 partial 的 store 早 ~35 µs 落地，整数校验和防读到半成品，kernel_dsv4_tp_flag_set_checked）；flag 折叠进载荷内核省一次 dispatch；RDMA 预热（Apple UC verbs 上向无 posted receive 的 QP 发送会永久杀死该方向，失败重建 QP/CQ）；prefill 32 MB 大门分级为 256×16 KB 链式 posted receive 窗口（~26 → 4-6 ms）；队列保活（M5 Max 空闲 3 s 后命令缓冲晚 600-800 ms 启动，每秒一个 1 KB 内核保温）；GPU 决定的负载均衡（按本 rank 专家 id 算共享专家车道互补切分，shift_q16，多领一个路由专家就让学生一半共享专家；计时驱动的 DS4_TP_LANE_BIAS 因输出依赖时序而默认关）；验证预投递 + 三态提交（ds4_tp_verify_commit_mode，TP 的 partial accept 也能直接提交前缀快照）。bbec200 后混合批 TP logits 与串行完全一致；DSpark/MTP 在 TP 下解禁但实测更慢（30-40 vs 47-48 t/s——每个验证行重流它的专家）。部署坑（不要 curl 等待中的 coordinator、恢复时双 rank 重建前缀）记入上游 docs/DISTRIBUTED.md。

Metal prefill 战役 + GLM SSD streaming 成熟化（融入 Part 6 §4p（新）/§4a、SSD Streaming 术语表） ​

Prefill 四抓手：MXFP4 路由专家上 TensorOps（kernel_mul_mm_id_mxfp4_*_mpp 复用 mul_mm_id 分组矩阵乘管线——数学相同、累加序不同，所以 exact 合同只签给 decode）；APE 位置表平铺（2048-token prefill 里 25K 个 dispatch 中 21.5K 个是在逐 token 编码位置表，折叠成一次平铺行拷贝）；路由激活复用/打包（跨专家批复用 gate/up 激活，两 rank prefill 打包 MXFP4 激活）；双缓冲 dense 内核 8 KiB threadgroup 修复（host 按 4 KiB 给——MTL_DEBUG_LAYER=1 的回归测试会抓住这类超配）。单 M5 prefill 790 t/s @2048 / 573 @16384 / 399 @65536。GLM streaming 主题是"缓存优先级可被运行期夺回"：内置热度加成封顶（decode 挤得掉冷种子）、从第一个路由层老化、运行期路由频率推翻静态 hotlist、≥256 token 续 chunk 前瞻一层读（8K append 63.2 → 90.7 t/s）、紧凑 prefill 保留 RoPE。实测 GLM Q4_K 121/104 t/s prefill、IQ2_XXS 16-token append 30.8 → 2.9 s；完整 GLM 走 server 仍会失速（已知回归 31d1992）。

DSpark：种子批处理、六 token 验证、EOS 契约（融入 Part 4 §5、DSpark 术语表） ​

种子批处理（be76012）：高置信种子融进验证器批量——被批种子是"正常选出的目标 token"，不改变输出分布；接受率差的窗口暂停起草 32 周期。Q2 C 哈希表 55.71 → 62.95 t/s；DGX Spark 三连跳到 31.41 t/s（散文 ~18.8——推测只对可预测文本回本）。六 token 验证（8d309ab）：种子 + 5 草稿需要前缀 1-4 各一份新鲜压缩器/索引器状态，spec_frontier_commit_prefix() 复用前缀快照槽；测试比较缓存行数而非文本，窗口管理器由无模型测试全量覆盖。EOS 契约（7a8c1b2）：草稿含内部 EOS 时在验证/提交前先截断——会话状态绝不跑到已返回 token 之外；7,217 用例 / 240,503 断言。CUDA 保留全部五个部分前缀、ROCm 保留四个。

ROCm gfx1151 第二轮（PR #992）（融入 Part 6 §4c） ​

版本匹配的 hipBLASLt prefill 方案：手调 solution 索引只在匹配的 hipBLASLt 版本上生效（g_hipblaslt_prefill_state 0/1/−1），版本漂移安全回退；DSpark 验证器注意力投影定宽方案（宽 2-6，5/6 行 ~2.55×）；分配前内存预算（ds4_rocm_memory.h 排除 CMA 页 + 2 GiB OS 余量，APU 识别，arena 按模型大小选——OOM 从运行中段提前成启动期拒绝）；死内核清除（−700 行）；跨后端 GLM 注意力掩码修复（dense 边界改由模型定义 + 同步归约存储，Strix Halo 续 prefill 21 → 68 t/s）。

服务器 / agent 工具加固 + 前缀条件化（融入 Part 5 §3/§3d/§3f（新）/§4、DSML 术语表、前缀条件化术语表（新）） ​

字面工具数据：工具体只转义自己的关闭分隔符（ds4_tool_text.h），agent/server 共用，KV 不因重渲染失效；GLM <tool_call> 方言接入同一套状态机，拆开的分隔符由采样跟踪器逐字节追踪；推测边界合同：工具/控制停止点丢弃未消费的推测 token、从恢复 logits 重采样，有效草稿跨边界保留，不修复截断动作、恢复受请求预算约束。会话回滚状态完整性（233eeb8）：GLM 递归 + 索引器状态随前缀回滚、TP 恢复双 rank 重建、live + 已计划会话都计入内存记账；test-session-state 锁定"多模态 API 的纯文本观察 == 普通消息 token"。对话前缀条件化（b0a147a）：--prefix-file 用 USER:/ASSISTANT: 前缀（assistant 补 EOS，GLM-DSA 除外）经 chat 模板渲染 + prefill 进 KV——上游引用 Many-shot Jailbreaking 说明为何必须显式 opt-in。agent：后台 bash 作业跨调用保活（bash_status/bash_stop，SIGKILL 截止期）、原子替换写入 + 并发修改检查、有界流式读取、终端重绘/UTF-8/Markdown 一揽子修复；make test-frontends 无模型跑全部前端回归。

ds4-eval hard 套件 + 文档重组（融入 Part 5 §5） ​

hard 套件（29aa726）：50 道题装进 ds4_eval_cases.c（1,118 行），四个有许可证来源（MMLU-Pro / OlympiadBench / LiveBench reasoning / NIST Juliet CC0）；答案判定支持有序序列与行集合；--suite {core,hard,hard-smoke} / --case-id / --validate-cases 等 CLI。质量工具链（PR #996）：validate_scores.py（评分器覆盖率）、compare_frontier_logits.py（全向量 logits 对比——长上下文回归从此有全向量参照）、GLM FP8 长上下文参考集 8 例。README 拆分为 12 篇 docs/ 指南（DISTRIBUTED/SPECULATIVE_DECODING/SSD_STREAMING/SERVER/PERFORMANCE/…），QA_BEFORE_RELEASES.md +707 行覆盖全部新主题的发布清单。

困惑与突破 ​

  • 困惑：DSpark 在单机上能把 decode 提到 62.95 t/s，为什么到了两台 Mac 的 TP 上反而比纯 decode 更慢（30-40 vs 47-48 t/s）？机器更多、算力更多，一次还验 6 个 token，怎么会负优化？
  • 突破：因为推测解码省的是"反复读权重的内存带宽"，而 TP 改变的是权重的存放方式。单机常驻模型里，一次 6 行验证批量共享同一份权重的读入，行越多越划算；TP + SSD streaming 下每行验证走各自不同的专家集合，6 行批量把"每 token 读一次半分片"的局容性打碎成"每行都要重流它的专家"（README 原话：each draft row re-streams its experts）——省下的带宽又在 SSD 重流里花回去了，还外加每层 RDMA 门的同步成本。同一份代码，单 M5 +8%、DGX Spark +60%、双 Mac TP 为负——加速技术没有普适，只有"在特定瓶颈下为真"；这也是为什么每个 DS4_TP_DISABLE_* / DS4_DSPARK_* 开关都要配 A/B 测量而不是拍脑袋默认。

上游同步 — 2026-09-11 ​

子模块：vendor/ds4 从 f62ca29 前进到 6289c51（main 分支，5 commits、~1.8K 行新增）。本轮是一轮"长任务生存能力"专题：agent 上下文压缩从"急救截断"进化为"任务不中断的常规换页"、多模态 live 上下文跨工具调用保持、教学提示模式 /hints 落地。已把内容融入各章节：

Agent 压缩即恢复（融入 Part 5 §4） ​

e33a5b2 把压缩变成生成中途可发生、发生后任务继续跑的机制：摘要空间预约（输入尺寸检查以压缩提示 + 摘要预算 ctx/8 + 64 为下限，超限消息干净拒绝、不挤坏当前对话）；满会话恢复（没有摘要空间就只摘要一个有界前缀、未摘要 token 全部逐字保留 retained_unsummarized）；tail 边界优先完整用户轮次（2×tail_budget 范围内向后找用户 token——"稍长的完整轮次"好过"丢了用户约束"）；未完成回复在同一条 assistant 消息内续写（保留 thinking 状态、结转输出预算 carried_generation，不插合成的"请继续"用户请求）；半截工具调用回滚重试（删残缺参数 + 系统说明"未执行，请用更小参数重试"）；图像块绝不切断（agent_compact_image_boundary）。摘要提示词纪律："不要解决未完成任务、记录实际发生了什么"。PTY 回归 test_agent_compaction.py 覆盖真模型压缩。

多模态 live 上下文保持（融入 Part 5 §7、视觉术语表） ​

81a7658（+844 行）把"图变了从零重建"细化为前缀匹配 + 图像位置重定基：ds4_session_vision_prefix_matches（历史图像跨度+指纹一致、新图在 live 前沿之后即可复用）、ds4_session_rebase_vision_state（live 续接经工具调用 ID 独立认证后，从检查点恢复图像 token 起点）、build_live_prompt_suffix（四条 live 续接路径统一：live token + 逐字新文本 tokenize + 新图像块拼接 + 终检）、visible_prompt_key（请求局部 marker nonce 归一化，字面文本伪装不成图像）、server_image_cache（32 项/128 MiB 字节精确键 LRU 嵌入缓存，同图不重跑 ViT）。图像工具结果三套 API 均可往返；M5 Max 上 GLM 5.3 Flash + Vision-Exp 验证 53K/56K 前缀。

教学提示 /hints（融入 Part 5 §4） ​

6546ae5（教学模式 PR #391 收敛版）+ 6289c51（终端徽章渲染）。运行期偏好不进会话元数据；完整指令每上下文一次、开关切换只发短说明；在安全 assistant 边界生效、不重建缓存；压缩/恢复会话重新确立策略（UNKNOWN 态先压过残留旧说明）。渲染是行首 > **Hint:** 的有界前缀保持状态机——白底青字徽章、无宽度依赖重绘。

其他 ​

74fbe54：QA_BEFORE_RELEASES.md +52 行——M5 Max 视觉客户端与压缩恢复的发布检查记录（聚焦 M5 Max 的测试结果与遗留的 DSpark 停止边界重建项）。

困惑与突破 ​

  • 困惑：压缩发生在生成的中间、且要在同一条 assistant 消息内续写——为什么要这么讲究？插一条合成的"请继续"用户消息，模型不就接着写了吗？
  • 突破：因为消息边界本身就是条件分布的一部分。对 next-token 预测来说，"assistant 消息 X 的后半段"和"用户请求继续之后的新回复"是两个完全不同的条件语境：后者会让模型把未完成的片段当作"已完成的事实"（跳过、复述、或干脆换个任务），thinking 块的开合状态也会错位。所以 resumed_start 直接指向同一条消息、<think> 扫描恢复思考态、输出预算结转——重建的不只是 token 序列，还有对话的语用结构。这与 Part 5 §3d 的"渲染文本前缀匹配"、DSML 的"字面工具数据"是同一条设计哲学：KV 可以重建，上下文的语义连续性必须原样保真。反例是半截工具调用——残缺的参数连"语用结构"都没法保真（执行它等于赌），所以宁可回滚重试。什么时候值得"原样续"、什么时候必须"推倒重来"，分界线就是能否验证。

上游同步 — 2026-09-12 ​

子模块：vendor/ds4 从 6289c51 前进到 bd66c40（main 分支，单提交 521 个文件——其中约 460 个是官方回复测试数据，核心改动 ~5K 行）。本轮是一整个新模型发布：DeepSeek V4.1 Flash（Metal 专用）——ds4 支持的第四个模型族，架构上最大的新意是一张 189 GiB、永远住在磁盘上的 Engram n-gram 记忆表。已把内容融入各章节：

DeepSeek V4.1 Flash + Engram 外部记忆（融入 Part 3 §8（新）、Engram 术语表（新）） ​

40 层 / 5120 维 / 384 专家取 6 / 候选块索引器 / 逐层 KV 压缩（前 20 层 2×）/ 原生 64K 上下文（1M 上限）。Engram 挂在第 1、14 层：4-gram 经词表压缩（NFKC→NFD→去重音→小写，129280→99092）与 multiply-xor 哈希，对 24 个素数桶（3 位置 × 8 头）取模得行号；行内 E4M3+E8M0 块缩放解码成 F32，由 engram_kv（[6144, 5·5120]）投影出 4 条超连接流各自的 key + 1 个共享 value，kernel_dsv41_engram_add 用缩放余弦门控加进残差（gate = sigmoid(copysign(√|cos|, cos))，BF16 对齐）。表永远不进内存：加载时 munmap，之后 F_NOCACHE + F_RDAHEAD=0 裸 pread——读过的页不许缓存，189 GiB 在常驻/streaming/TP 任何模式下都住在 SSD 上。哈希状态只有 3 个 int（复制即快照），快照恢复重放最后 3 个 token 即可重建。转换工具链 deepseek41_*.py 四件套逐位复现官方 inference/engram.py（素数表 sympy nextprime、乘子固定种子 RNG——C 侧只消费数组）。分词器也换了血：<｜System｜> 和 <tool_result> 成为真特殊 token。

Metal V4.1 执行引擎（融入 Part 6 §4q（新）/§4e/§4j、SSD streaming 术语表） ​

全新图路径 ds41_gpu_graph + 新内核源 metal/dsv41.metal。三个抓手：层主序宽批 prefill（≤8192 token，算第 il 层时预读 il+1 层，≥16384 token 再 mlock 前 20 层专家；~3 GiB 紧凑 carry 换页）；原生会话解码 2–8 路（逐位对齐标量——复用 §4n tiny-pair 内核、projection_rows 保持逐行算术；TP 3–8 路、SSD streaming 与图像会话走有序兜底）；数值契约（decode 签 exact 合同，层主序 prefill 的分组 GEMM 改归约顺序——Q4 batched prefill 概率分数略降是已知已记录的差异，top-token 总数不变但上游明说不算等价证明）。QA §17 的教训：Q4 首跑把一台 M5 压到重启（PRO/Q4 间接专家表整层驻留、123.4 GiB wired vs 计划 76.69）→ 修成选中专家缓存；8 行 decode 超 2e-4 容差 → 标量专家归约修回；512 token 热 SSD 批处理阈值因长 NLL 回归被否。实测 512 GiB M3 Ultra Q4：prefill 342–716 t/s、decode ~18.1–18.7 t/s。TP 协议 10→12：同步检查点帧（取消两 rank 协商一致）、RDMA 全发送 SIGNALED + 窗口屏障修 UC 竞态。

think-level / 模板 / agent（融入 Part 5 §3g（新）/§4/§6） ​

V4.1 把 thinking 变连续旋钮：--think-level 1..100（--think=75、--think-max=100、0 关闭），编码进 ds4_think_mode 枚举（基值 1000 + 档位），渲染成紧跟 BOS 的 "Reasoning Effort: N" 系统文本——换档即改最前 token、KV 整段重建；agent 恢复会话时逐档渲染 101 个候选前缀按最长匹配还原。服务器侧：合并轮次模板 render_deepseek41_chat_ordered（空 think 块恒定发出、完整 prompt 与 live 尾部同一渲染器）、DSML 空格拼写变体、并行工具结果按调用顺序重排（ds41_order_messages）。agent 的 bash 工具改 posix_spawn——推理进程里躺着几百 GiB wired Metal 映射，fork 复制页表本身就能压垮内存。

V4.1 视觉（融入 Part 5 §7、视觉术语表） ​

第三条视觉管线：共享 patch 几何与 ViT 骨架（ds4_image_preprocess_deepseek 带 v41 开关），但 token 预算 1024（V4 Flash 是 384）、无 8:1 长宽比钳制、朴素行主序 token 流（不再行对交织）。编码器 266 个 BF16 张量逐字节直拷，投影 5120 维。图像位置同时抑制 Engram；含图会话排除在原生批处理外。

困惑与突破 ​

  • 困惑：为什么 Engram 的哈希不处理冲突？教科书哈希表都要探测/链式兜底，这里碰撞了不就是读错行、注错记忆吗？
  • 突破：因为它是学习出来的冗余，不是工程冗余。同一个 4-gram 探 3 个位置 × 8 个头共 24 行，投影后还有余弦门控——训练时表内容、乘子、素数、投影、门控是一起学的：碰撞行如果对任务有害，梯度会把它的 value 学成"门控天然压小"的样子。工程哈希表的冲突处理保证"任意键都不读错"；Engram 只要"期望上读对、错了不致命"——它是在浮点噪声里训练出来的，天生对一行误差鲁棒。这和 KDA"固定状态必有损、靠混合架构兜底"是同一种设计品味：不追求机制上无损，追求端到端可学。另一层反差在 I/O：常规权重靠 mmap 缺页缓存（越读越快），Engram 反过来 F_NOCACHE 禁止缓存——189 GiB 若走 page cache 会把整台机器的内存策略挤垮，而每 token 12.7 KiB 的随机读 SSD 本来就扛得住。"什么该缓存"本身就是架构决策。

上游同步 — 2026-09-20 ​

子模块：vendor/ds4 从 bd66c40 前进到 0aaea5a（main 分支，79 commits；核心改动 ~22K 行，其余是官方续写测试数据与 Qwen 质量夹具）。本轮两件大事：第五个模型族 Qwen3.8 Flash Next（GDN 混合架构 + 盘上 n-gram PLE + Metal/CUDA 双后端 + 原生批处理与批处理推测解码），以及 V4.1 Flash 上 CUDA（单 Spark SSD streaming + 双 Spark RoCE TP）。另有服务器槽位路由重构与一串流式/渲染修复。已把内容融入各章节：

Qwen3.8 Flash Next —— 第五个模型族（融入 Part 3 §9（新）、GDN 术语表（新）） ​

48 trunk 层（3 GDN + 1 门控 GQA，interval=4）+ 1 MTP 块 / 2560 维 / 词表 248320。三条已知路线的合流：GDN（qwen4_ref_linear）是 KDA 近亲——卷积→衰减 exp(A·softplus(a+dt))→L2 归一化→delta 规则 (v−H·k)·β 残差写入→sigmoid 输出门，48 v 头配 16 k 头，3/4 的层零 KV；门控 GQA 是 24Q:2KV（12:1）+ Q 投影半宽 sigmoid 门 + QSA 索引器（4-token 块、relu 求和、top-512 块 = 2048 token）+ 64/256 维部分 RoPE、3D 位置表；PLE 是 Engram 表亲——第 1 层挂 95.37 GiB 原生 BF16 3-gram 表（multiply-xor 哈希 16 行/token、pread 直读、拒绝量化过的表）、同款 sigmoid(copysign(√|g|,g)) 缩放余弦门控 + 膨胀深度卷积。MoE 512 取 10 + 1 共享（共享专家带学习到的 sigmoid 缩放）；4 流超连接（低秩 320）。Q2 工件 41.73 GiB 常驻（gate/up imatrix IQ2_XXS、down Q2_K 640→768 列填充），64 GB Mac 起步；原生 262144 上下文（DS4_QWEN4_YARN_FACTOR）。视觉直接消费 llama.cpp 的 Qwen3-VL mmproj（第四族视觉管线）；qwen35 预分词的 Unicode 类由 Python regex 模块生成 ds4_qwen4_unicode.inc。

Qwen GPU 图与批处理战役（融入 Part 6 §4s（新）、Part 5 §6、会话批处理术语表） ​

GPU 图一条图服务 decode 与批量 prefill（f32 瞬态按 cap_tokens 定尺寸）；内存契约：瞬态可共享（同一时刻只跑一个 forward）、递归状态/缓存/PLE 历史/快照每会话私有——所以批处理必须是"多行一次 forward"。原生跨会话批处理把 16 session = 16 倍权重流量变成"dense 工作整批一次 + 行表寻址的 GDN/注意力/PLE 逐会话状态"（~2,300 个短 dispatch → 表寻址批量内核），会话瞬态并共享 arena（每 session 内存 ÷3.6）；C=16 @M5 Max：50 → 143 t/s 聚合。批处理推测解码：token 作行 0、草稿搭车作行 1 整批验证，预测器层整批跑一遍，贪心接受、失败恢复行快照、预测器失败整批失效；逐周期"整批同进退"的起草决策 (1+p)·c_plain > c_spec（EMA 定价，稳态周期才计价，批宽变化重置）。盘上 n-gram 并发读（≥2 行走 dispatch_apply_f）：C=1 普通 51.3→54.5、MTP 代码 69.9→78.9 t/s，logits 逐位相同。MTP 单会话侧：默认深度 1、"一级接受完美"时链式二级草稿（3 行验证），DS4_QWEN4_MTP_DEPTH 固定；工具采样边界回卷恢复 pre-verify 递归快照而非重放图（融入 Part 4 §5）。

V4.1 CUDA + 双 Spark TP（融入 Part 6 §4r（新）/§4b/§4j、SSD streaming 术语表、张量并行术语表） ​

同一个模型两种存储哲学：Metal 是共享 mmap 缺页即算，CUDA 是有界 device 专家缓存 + 恰好两个 8 MiB pinned staging 缓冲 + 一个读线程预读下一层（≥2K token 才开，预留槽拷贝全部完成才发布）；中等追加"一次层扫"（每层只拜访一次）3.2K append ~55→98 t/s；Q2 SSD serving 最多 8 行 decode 批。实测：自动缓存 3.2K prefill 49-54→93-96 t/s、64 GiB 预算 32K prefill 384 t/s，decode 微降 5% 换整段请求 87→59s。双 Spark TP：Linux RoCEv2 RC 队列对 + pinned host 中转（RoCE 非 GPUDirect）、同步交换模型（对照 Metal 服务线程异步门）、协议 → 14（所有 peer 一起升级）；单会话 21.9 t/s、8 会话聚合 ~28（5 个起批）、32K prefill ~400 t/s。decode 图捕获每 rank 120 区域。配套：会话并发基准（vLLM 风格 TTFT/ITL/TPOT + p50-p99，闭合循环与 Poisson 两种到达；按引擎内存记账定价网格，不完整的流整份拒绝）。

服务器槽位路由重构（融入 Part 5 §6） ​

两步：4d7028e 把每条 RAM 复用路径拆成"廉价探测 + 物化"，阶梯（responses-visible → tool-output ×2 → GLM rewind → token → thinking-visible → text）只定义一次、执行与评分消费同一探测——修掉"平行 subagent 被吸到主 agent 长上下文槽上零收益驱逐"的事故；分数从检查点全长改成实际可复用前缀长。ada805a 给零复用受害者分陈旧度层：空槽 > 闲置超 1h（最闲置优先、不看长度——已完成的长对话先回收）> 活跃（-live 最短优先）；陈旧只降受害者优先级、从不阻止复用；驱逐派发时记日志。5994f06 给 token/rewind 层加历史图像位置精确校验（token 相等证明不了图没挪位）。

流式/渲染修复（融入 Part 5 §3h（新）/§7、视觉术语表） ​

Qwen 工具内容流式（guard_second_reasoning 是 DeepSeek 行为的补丁，按模型语法豁免 Qwen）、跨 chunk 分裂的 </think> 标记、参数值里的字面 </tool_call>（服务器用结构化定界符搜索、agent 状态机只等自己的定界符）；GLM 工具块内联渲染（重建路径多俩换行 → 检查点 key 差 3 字节 → 重启重 prefill 3800/11992 token）；reasoning_effort 接受数值字符串（pi.dev 兼容）；/v1/models 报实际加载的 GLM id；Iris 渐进 JPEG ZRL 修复（run=15 被当 run=16、吞掉下一符号的修正 bit——位流失步，与 djpeg 逐位对齐修复）；GLM Metal 填充选择掩码（selected_rows_valid 的承诺由生产者证明而非调用者许愿）；V4.1 decode 用上并发 Engram reader（+9.9%，输出逐字节相同）与路由 softplus log1p 级数（融入 Part 6 §4q/§4l）。

困惑与突破 ​

  • 困惑：批处理推测解码为什么要"整批同进退"——每个 session 按自己的草稿置信度独立开关不是更精细吗？
  • 突破：因为行成本不是线性的。上游的测量：16 行 80 ms、23 行 126 ms、32 行 149 ms——第 17 行会打开每个 dense 投影的第二个 tile，跨过这道台阶之后行几乎免费。在这个成本结构下，"省掉一个 session 的草稿行"省不到钱（它多半在免费区间里），"多开一个草稿行"也多半不额外花钱——单 session 粒度的开关在错误的尺度上优化。正确的决策粒度必须对齐固定成本的台阶（tile 边界），于是退一级到整批同进退，用 (1+p)·c_plain > c_spec 的 EMA 定价在周期尺度上决策。这与 §4j"两路 TP 摊不平批量交换的固定开销就退有序兜底"、DSpark"种子 token 融进验证器批量"是同一条经济学：存在台阶状固定成本时，调度粒度要对齐台阶，而不是对齐请求。反例恰好也在本轮：盘上 n-gram 的串行读（每 token 16 行 × 320 B，1.3 ms）就是线性成本——它的解法不是改调度粒度，而是并发扇出（读快 5 倍）。先判断成本曲线是台阶还是直线，再选工具。

出链