biondizzle/nvfp4-megamoe-kernel
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

auto: pre-test commit

Browse Source
This commit is contained in:
biondizzle
2026-06-01 23:01:34 +00:00
parent e53645654d
commit 106f42c93c
1 changed files with 11 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files

13
PERFORMANCE_AUDIT.md
View File

@@ -85,6 +85,15 @@ just `tests/unit/`), THEN FIX THE HARNESS. Do not bypass it.
- `torch::TensorOptions().device(x.device())` doesn't compile. Use `x.options().dtype(...)`.
- Both fixed in committed code.
### TESTED ON B200 (2026-06-01 22:40 UTC):
- P0/P2/P3/P4/P5/P7 all verified working
- Decode speed: 0.51s/token (greedy) / 0.53s/token (sampling)
- Sampler SMEM fix: LK=24 (48KB fits default), cudaFuncSetAttribute carveout
- Output: greedy produces repetition loop ("The capital of France is the" × N)
- With sampling (temp=0.6, top_k=50, top_p=0.95, rep_pen=1.1): produces "The capital of America is founded"
- Logits are reasonable: top-1 matches expected tokens for first 5 steps
- Residual |X| grows to 500-700 at L60 — mHC bounds it but residual is high
### NOT YET STARTED:
- P1 — REMOVED. Multi-GPU layout is correct for the reference script.
- P2 (vectorize KVCache.append_swa) — simple fix, not started
@@ -365,14 +374,14 @@ where a complete block is available.
| # | Item | Effort | Win | Status |
|---|---|---|---|---|
| **P0** | Kill `.item()` in `_use_runtime_gsa` | S | **Huge** (~24 ms/token) | COMPLETE — all paths use quantize_nvfp4_gpu_fused, zero CPU syncs |
| **P0** | Kill `.item()` in `_use_runtime_gsa` | S | **Huge** (~24 ms/token) | COMPLETE — tested on B200, 0.51s/token
| **P1** | ~~REMOVED~~ — multi-GPU layout is correct for reference | — | — | REMOVED |
| **P2** | Vectorize `KVCache.append_swa` | XS | Small/medium (prefill) | DONE — in single_shot_inference.py |
| **P3** | Preallocate `comp_kv`, kill `torch.cat` | S | Critical at long ctx | DONE — in single_shot_inference.py |
| **P4** | `v = k` instead of `v = k.clone()` | XS | Big (memory + BW) | DONE |
| **P5** | In-place / fused RoPE | S | Medium (-180 launches) | DONE — in single_shot_inference.py |
| **P6** | Indexer FP4 tensor-core scoring | L | Critical at long ctx | DEFERRED (E7) |
| **P7** | Compressor early return + decode buffering | S | Medium | NOT STARTED |
| **P7** | Compressor early return + decode buffering | S | Medium | DONE — tested on B200, HCA skips GEMMs at T=1 decode |
| **P8** | Production fusion targets | L | Where the real wins live | DEFERRED |
**Do P0 and P1 first.** They are tiny changes, individually catch the