docs: record sm90 megamoe phase5 dev log

MEGAMOE_SM90_DEV.md

@@ -166,3 +166,53 @@
 - 后续待办：
   - Phase 5：实现 L2 GEMM consumer，消费 `l2_arrival_mask` 和 Phase 4 产出的 FP8 intermediate activations。
   - 后续优化：将 Phase 4 global store 替换为 STSM + TMA store，并扩展 correctness 覆盖多 expert、非首 pool block 和更多 shapes。
+
+## 2026-06-18 15:17:31 CST
+
+- 对应 clean 工作 commit：`9bd05196051dd3f87ebc1eb2ec0c417160ed0418`
+- 用户请求摘要：阅读 Phase 5 计划并执行，实现 SM90 MegaMoE L2 GEMM + scatter 的首个可验证版本。
+- 本次提交的核心改动：
+  - 扩展 SM90 symmetric buffer views，新增 `combine_acts: bf16[num_topk, num_max_tokens_per_rank, hidden]`，用于 Phase 5 验证 L2 scatter 输出。
+  - SM90 kernel TMA producer 支持 `SM90BlockPhase::Linear2`：等待 `l2_arrival_mask == ((1 << (2 * num_k_blocks)) - 1)` 后加载 `l2_acts` 和 `l2_weights`。
+  - SM90 math path 支持 L2 WGMMA：对每个 K=128 block 拆分两个 per-64 half，分别累积 raw FP8 WGMMA，再应用 `l2_acts_sf` 和 `l2_weights_sf` 得到 FP32 L2 output。
+  - 新增 BF16 direct scatter：根据 `token_src_metadata` 将 L2 output 写入本地/远端 `combine_acts` 对应 top-k slot 和 token 位置。
+  - 新增 Phase 5 correctness 脚本 `megamoe_dev_test_scripts/phase5/l2_gemm_scatter_correctness.py`，验证 single-rank 多 expert 下 `combine_acts` 与 PyTorch L2 GEMM BF16 reference 一致。
+  - 更新 Phase 1 interface smoke，检查 `combine_acts` shape 和 dtype。
+- 关键文件列表：
+  - `csrc/apis/sm90_mega.hpp`
+  - `deep_gemm/include/deep_gemm/impls/sm90_fp8_mega_moe.cuh`
+  - `deep_gemm/mega/__init__.py`
+  - `megamoe_dev_test_scripts/phase1/interface_smoke.py`
+  - `megamoe_dev_test_scripts/phase5/l2_gemm_scatter_correctness.py`
+- 详细开发流程：
+  - 本地阅读 `MEGAMOE_SM90_DESIGN.md` Phase 5 计划，确认目标为 L2 WGMMA pipeline、per-64 dual-half SF、BF16 epilogue 和 scatter 到 combine buffer。
+  - 本地对照 SM100 MegaMoE 的 Linear2 TMA wait/load、L2 BF16 epilogue 和 combine buffer scatter 逻辑，确认 SM90 首版采用 direct global scatter，先不实现 STSM staging。
+  - 本地扩展 SM90 host buffer slicing 和 Python `SymmBuffer`，暴露 `combine_acts`。
+  - 本地在 SM90 kernel 中重建 `combine_token_buffer` layout，并在 TMA producer 中为 Linear2 增加 `l2_arrival_mask` wait 和 L2 acts/weights TMA load。
+  - 本地在 math branch 中保留 Linear1 L1 epilogue，同时新增 Linear2 dual-half WGMMA/scaling/BF16 scatter 分支。
+  - 本地新增 `megamoe_dev_test_scripts/phase5/l2_gemm_scatter_correctness.py`，固定 `num_tokens=256`、`num_experts=2`、`num_topk=1`，让两个 expert 各接收 128 tokens，覆盖 multi-expert single-rank Phase 5 milestone。
+  - 本地静态检查：`python3 -m py_compile megamoe_dev_test_scripts/phase5/l2_gemm_scatter_correctness.py megamoe_dev_test_scripts/phase4/l1_epilogue_correctness.py megamoe_dev_test_scripts/phase1/interface_smoke.py`，通过。
+  - 本地静态检查：`git diff --check`，通过。
+  - 远端同步：`rsync -avz --relative csrc/apis/sm90_mega.hpp deep_gemm/include/deep_gemm/impls/sm90_fp8_mega_moe.cuh deep_gemm/mega/__init__.py megamoe_dev_test_scripts/phase1/interface_smoke.py megamoe_dev_test_scripts/phase5/l2_gemm_scatter_correctness.py g0034:/mnt/beegfs/lxy/DeepGEMM/`。
+  - 远端 build：`ssh g0034 "docker exec sglang-lxy bash -lc 'cd /mnt/beegfs/lxy/DeepGEMM && timeout 60s bash develop.sh'"`，通过。
+  - 远端 Phase 5 single-rank 首次 correctness 触发 JIT 后失败在测试 reference：脚本直接使用 `expert_recv_count_sum`，未 mask 低 32 位 token count；本地修复为 `count = value & 0xffffffff` 后重新同步测试脚本。
+  - 远端 Phase 5 single-rank correctness：`ssh g0034 "docker exec sglang-lxy bash -lc 'cd /mnt/beegfs/lxy/DeepGEMM && timeout 60s python3 megamoe_dev_test_scripts/phase5/l2_gemm_scatter_correctness.py'"`，通过。
+  - 远端 Phase 1 single-rank/2-rank regression：`interface_smoke.py`，通过。
+  - 远端 Phase 2 single-rank/2-rank regression：`dispatch_only_correctness.py --tokens-list 192`，通过。
+  - 远端 Phase 3 single-rank/2-rank regression：`l1_wgmma_single_tile.py`，通过，`max_diff=0.006541`。
+  - 远端 Phase 4 single-rank/2-rank regression：`l1_epilogue_correctness.py`，通过。
+  - 远端 Phase 4 8-rank regression：`ssh g0034 "docker exec sglang-lxy bash -lc 'cd /mnt/beegfs/lxy/DeepGEMM && CUDA_VISIBLE_DEVICES=0,1,2,3,4,5,6,7 timeout 60s torchrun --standalone --nproc_per_node=8 megamoe_dev_test_scripts/phase4/l1_epilogue_correctness.py'"`，通过。
+  - 提交前本地检查：`python3 -m py_compile ...`、`git diff --check`、`git diff --staged --check`，均通过；检查 touched files 未发现 `// DEBUG` 或 `# DEBUG` debug 代码。
+- 测试结果和已知问题：
+  - Phase 5 build 通过。
+  - Phase 5 single-rank multi-expert L2 GEMM scatter correctness 通过。
+  - Phase 1/2/3/4 single-rank 和 2-rank regressions 通过。
+  - Phase 4 8-rank regression 通过，说明新增 Linear2 路径未导致现有 multi-rank Phase 4 correctness hang/crash。
+  - single-rank 仍出现 PyTorch symmetric memory warning：`cuMulticastCreate failed with: "invalid argument"`，当前不影响 correctness。
+  - Phase 5 当前采用 direct global BF16 scatter，没有实现计划中的 BF16 STSM 到 `smem_cd_l2` 再 shared-load scatter。
+  - Phase 5 correctness 当前覆盖 single-rank 多 expert；multi-rank L2 scatter 数值 reference、top-k>1、非首 pool block、多 K block/更多 shape 尚未覆盖。
+  - Phase 6 combine reduce 和最终 `y` 写回尚未实现，因此端到端 fused output 仍未完成。
+- 后续待办：
+  - Phase 6：实现 combine reduce、NVLink barrier、workspace cleanup 和最终 `y` 写回。
+  - 扩展 Phase 5 测试到 multi-rank scatter reference、top-k>1、更多 hidden/intermediate shape 和非首 pool block。
+  - 后续性能优化：将 direct global scatter 优化为 BF16 STSM/shared-load scatter，降低 store 指令和改善 coalescing。