docs: fill the Phase 19 gap, refresh README/roadmap to actual state
- docs/19-gpt-oss-moe.md: the numbered series jumped 18->20; write up gpt-oss arch deltas, harmony pitfalls, and the two CUDA debugging postmortems (fully-masked-tile NaN in flash-attention sinks; pre-__syncthreads early return reading uninitialized smem in the decode GEMV) — the highest-value learning content of that phase. - README: models/perf/capabilities were frozen at the Qwen3-only era; now lists gpt-oss MoE, TP/PP, FP8/MXFP4, sparse MoE, and the llama.cpp standing. - Roadmap: record where reality diverged from the plan at Phase 18+, add milestone entries and the ranked next-phase candidates (21 CUDA-graph MoE decode, 22 non-expert quant, 23 sparse prefill). - sparse-moe benchmark doc: post-review-fix numbers. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
This commit is contained in:
118
docs/19-gpt-oss-moe.md
Normal file
118
docs/19-gpt-oss-moe.md
Normal file
@@ -0,0 +1,118 @@
|
||||
# Phase 19: gpt-oss-20b — MoE 模型支持与两次 CUDA 调试实战
|
||||
|
||||
> 目标:支持 OpenAI gpt-oss-20b(32 专家 top-4 MoE),GSM8K 精度对齐 llama.cpp,
|
||||
> 并以此为载体做 FP8 / MXFP4 量化。本文档事后整理,重点放在**踩过的坑**:
|
||||
> 两个教科书级的 CUDA bug 排查过程比结论本身更有学习价值。
|
||||
>
|
||||
> 后续:`docs/20-sparse-moe.md`(稀疏化),benchmark 数据见
|
||||
> `docs/benchmarks/{fp8-quantization,mxfp4-and-llama-decode,sparse-moe}.md`。
|
||||
|
||||
## 1. 模型架构(与 Qwen3 的差异点)
|
||||
|
||||
gpt-oss-20b(`config.json`,已在 dash5 验证):
|
||||
|
||||
| 项 | 值 | 说明 |
|
||||
|---|---|---|
|
||||
| layers / hidden | 24 / 2880 | hidden **不是** 128 的倍数的来源(2880 = 22.5×128) |
|
||||
| heads | 64 Q / 8 KV,head_dim **64** | head_dim ≠ hidden/heads(64×64=4096>2880),GQA n_rep=8 |
|
||||
| MoE | 32 experts,top-4,expert inter 2880 | router 是普通 Linear [2880→32] + bias |
|
||||
| attention | **交替 sliding(128)/full**,layer 0 是 sliding | 每层带 **attention sinks**(每 head 一个可学习标量) |
|
||||
| RoPE | YaRN(theta 150000, factor 32, orig 4096) | attn_factor = 0.1·ln(32)+1 乘在 cos/sin 上 |
|
||||
| 激活 | clamp 后的 GLU | gate=gu[::2], up=gu[1::2](**交错**), gate≤7, up∈[-7,7], glu=gate·σ(1.702·gate), h=(up+1)·glu |
|
||||
| 词表 | 201088 | EOS 是**列表** [200002,199999,200012] = `<|return|>`/`<|endoftext|>`/`<|call|>` |
|
||||
| 其它 | attention_bias=true | q/k/v/o 全部带 bias(Qwen3 没有) |
|
||||
|
||||
**Harmony 对话格式**:gpt-oss 不是普通 chat template,输出分 channel
|
||||
(`analysis`=思维链,`final`=正式回答),控制 token `<|start|>/<|channel|>/<|message|>/<|end|>`。
|
||||
三个坑:(1) system 消息必须含 `Reasoning:` 等 canonical 行,缺了模型 OOD、
|
||||
channel 选择不稳定;(2) repetition penalty 会惩罚必须重复出现的控制 token,
|
||||
导致模型只输出 analysis 不出 final(MoE 默认关掉);(3) 服务端要用多 EOS 判停。
|
||||
|
||||
## 2. MoE 前向(dense 版,Phase 20 之前)
|
||||
|
||||
```text
|
||||
router GEMV → topk_softmax(GPU)→ moe_replicate(复制到全部本地专家)
|
||||
→ batched GEMM gate_up → bias → GLU → batched GEMM down → bias
|
||||
→ weighted_sum(只取 top-4)→ all-reduce
|
||||
```
|
||||
|
||||
要点:top-k 的专家编号始终留在 GPU(`topk_ids`),host 不同步;
|
||||
dense 的代价(每 token 读全部专家权重)在 Phase 20 用 sparse GEMV 解决。
|
||||
TP 用 **expert parallelism**:rank r 拥有专家 [r·E/world, (r+1)·E/world),
|
||||
weighted_sum 里按 `expert_start + local_experts` 过滤非本地命中,
|
||||
all-reduce 把各 rank 的部分和加起来——这要求"跳过"语义而不是"乘 0"。
|
||||
|
||||
## 3. CUDA 调试实战 ①:prefill NaN(flash-attention sinks)
|
||||
|
||||
**症状**:长 prompt(≳192 token)prefill 后输出全 NaN → argmax 落在
|
||||
token 201087(`max_by` 平局取最后)或 token 0(`!`)。短 prompt 完全正常。
|
||||
|
||||
**定位手法**:给每个 stage 加 NaN 检查(环境变量开关,事后移除),
|
||||
二分出第一个出 NaN 的位置:layer-0 的 `flash_attention_sinks` 输出,
|
||||
而它的 q/k 输入是干净的 → bug 在 kernel 内部。
|
||||
|
||||
**根因**:causal 跳过逻辑只剔除"完全在未来"的 kv tile;一个完全滑出
|
||||
sliding window(128)的**过去** tile 仍被处理,所有 key 都被 mask 成 -inf
|
||||
→ `row_max = -inf` → online softmax 里 `expf(-inf-(-inf)) = NaN`,
|
||||
下一个有效 tile 的修正项 `0·NaN` 把整行毒掉。
|
||||
|
||||
**修复**:`row_max == -INFINITY` 的 tile 直接跳过(贡献为零)。
|
||||
**教训**:online softmax 的"空 tile"是边界条件标配;decode kernel 早就
|
||||
防了这个(`local_max==-INFINITY` guard),prefill kernel 漏了——
|
||||
**同一逻辑的两份实现要做同样的边界测试**。触发阈值 ~192 token 解释了
|
||||
"短测试全过、长对话必炸"的诡异表象。
|
||||
|
||||
## 4. CUDA 调试实战 ②:decode 间歇性乱码(GEMV 未初始化共享内存)
|
||||
|
||||
**症状**:同一 prompt ~70% 的运行在第二轮对话或长生成中突然输出
|
||||
`!!!!`/token 201087/NaN logits,**间歇性** → 不是确定性逻辑错误,
|
||||
是竞态或未初始化读。只有 gpt-oss 出问题,Qwen3 从不复现。
|
||||
|
||||
**定位**:逐 stage 检查,第一个出问题的是 decode 的 o_proj 输出
|
||||
(maxabs≈1e33),输入干净 → M=1 的 GEMV kernel。
|
||||
|
||||
**根因**(`gemv.cu`):
|
||||
|
||||
```cuda
|
||||
if (col >= N) return; // ← 在协作加载 x_shared 和 __syncthreads 之前!
|
||||
...cooperative load + __syncthreads()...
|
||||
```
|
||||
|
||||
当 `N % 128 != 0` 时,最后一个 block 的越界线程提前退出,**没参与**
|
||||
共享内存装载;在界线程读到未初始化的 smem(且 `__syncthreads` 在有线程
|
||||
已退出时是 UB)。命中条件:n=2880 的矩阵(o_proj、MoE gate_up/down)——
|
||||
2880 % 128 ≠ 0;而 Qwen3 所有维度都是 128 对齐的,**所以"只有 gpt-oss
|
||||
不稳定"**。q/k/v(4096)、lm_head(201088)对齐,幸免。
|
||||
|
||||
**修复**:所有线程先完成装载 + barrier,`col >= N` 检查移到 syncthreads
|
||||
**之后**。
|
||||
|
||||
**教训**:`__syncthreads()` 之前的任何 early-return 必须是 **block-uniform**
|
||||
的。Phase 20 的 sparse GEMV 专门遵守了这条(整个 block 基于同一个
|
||||
`topk_ids` 值统一退出,发生在装载之前)。
|
||||
|
||||
**修复后的验证**:GSM8K 全量 1319 题,xserv 94.5% vs llama.cpp 94.4%
|
||||
——统计上同一水平,证明两个 kernel bug 就是之前 55% vs 95% 差距的全部原因。
|
||||
|
||||
## 5. 量化(详见 benchmark 文档)
|
||||
|
||||
- **FP8 W8A8**(`tools/quantize_fp8.py`):per-expert 标量 scale,权重转置
|
||||
存 [E,N,K] 喂 cuBLASLt(Blackwell 要求 transA=T)。两个性能坑:
|
||||
(1) 每次调用重建 plan + 跑 heuristic → 比 BF16 还慢,修复 = per-shape
|
||||
plan cache;(2) 逐专家发射 ~768 个小 GEMM,修复 = 单条 strided-batched
|
||||
调用 + 把 scale 移到融合的 post-scale kernel。最终 1.41× vs BF16。
|
||||
- **MXFP4 W4A16**(`tools/quantize_mxfp4.py`):E2M1 + per-32 UE8M0 块 scale,
|
||||
13GB 模型,贪心输出与 BF16 逐字一致,但手写 dequant-GEMV 打不过
|
||||
cuBLASLt FP8(带宽效率差),定位为省显存方案。
|
||||
- 检测方式:safetensors 的 dtype/scale 秩自动识别,loader 无需配置。
|
||||
|
||||
## 6. 本阶段的工具沉淀
|
||||
|
||||
- `bench-gpt-oss`:in-process 推理 + `--forced`(teacher-forced prefill
|
||||
top-1)/`--forced-decode`(沿参考轨迹逐位置 top-1)——分离"前向算错"
|
||||
和"贪心轨迹分叉"的利器。
|
||||
- `tools/eval_gsm8k_fast.py`(持久 xserv-chat 管道)、
|
||||
`tools/xserv_vs_llama.py`(warm-server 同机对打,计入 llama 的
|
||||
reasoning_content)。
|
||||
- 经验:**贪心解码不是逐位可复现的**(cuBLAS 非确定性会翻转后段 argmax),
|
||||
多卡正确性要用"单卡×2 + 多卡×2 互相比",精度要用基准集而不是逐字 diff。
|
||||
Reference in New Issue
Block a user