Skip to content

学习日志


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 个)
  • 关键洞察:专用引擎 > 通用框架——固定架构换来简化和优化

困惑与突破

  • 困惑:为什么 makeds4_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 + memsetcalloc 更安全?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.mdmisc/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/ds4c463029 前进到 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)用 hipccds4_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/ds4e34a808 前进到 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)题里行首的 II think...)、AA 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 选项的有效范围只到 DI 根本不在有效字母范围内,find_answer_letterc >= 'A' && c <= max_answer 直接跳过它——所以无害。只有当干扰字母本身落在 [A, max_answer] 区间内(10 选项 A–J 时 I 进入了范围),它才有机会"假装是答案"。这是"边界条件 + 数据集规模"共同决定的可触达性:bug 一直存在,但只有 24 道 10 选项用例能触发它。

上游同步 — 2026-07-29

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

GLM 5.2 —— 第二种模型架构(融入 Part 3 §5Part 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.hds4_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/ds454b36ed 前进到 b030961(main 分支,86 commits)。本轮是 Flash 0731(DS4F)落地 + CUDA 性能大跃进:新增 MXFP4 量化格式(CUDA/Metal 原生推理 + 无损 GGUF 转换)、DeepSeek V4 Flash 0731 检查点成为 L40S 默认、把 batched-serving fork 的 vendored MMQdecode-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-benchPromessi 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 tier0f62c62):从 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 回退。配套 6f087f4make 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_120ab1f8893):RTX 50 / PRO Blackwell 编为 compute_120a/sm_120a,这样才接受 block-scaled MXFP4 指令;兼容 sm_120sm_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.7a51e6ec)——放更多后缀过验证,换更高接受率。
  • 贪心同一性修复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 状态)才安全。