Skip to content

Commit 9719695

Browse files
apstenku123apstenku123
committed
lever5(mamba3-compare): gb10 per-region ours-vs-cppmega bench + doc skeleton
scratch/mamba3_m2rnn_compare.py times, at the SAME config, OUR path_c gridded CUDA SSD per-region cost (F0/F1/F2 fwd, B0/B1/B2 bwd + host-staging chain delta) vs cppmega's SAME-model ops (mamba_chunk_scan_combined Megatron SSD Triton fwd+bwd, m2rnn_scan_triton fused Triton fwd+bwd), to LOCATE where path_c loses time on the identical model. Prints a machine-parseable RESULT json (per-region ms, ours/cppmega ratios, fwd/bwd host-staging fraction = the §17 launch-overhead term). RULE #1 honest attribution: ours-M2RNN has NO CUDA twin (Metal-only Path-C) -> reported NOT-RUNNABLE-ON-CUDA, never fabricated; missing mamba_ssm/triton -> hard error surfaced, not swallowed. Kernel-call ABIs reuse probe_chunked_{scan,backward}_cuda_gb10.py verbatim (one source of truth). docs/MAMBA3-PATHC-VS-CPPMEGA.md = methodology + pending-numbers table. local-check: py_compile OK, ast OK, cross-symbol OK (all 11 imported syms present), Mac dry-run exits 2 (no CUDA, expected). gb10 phase fills the MEASURED columns.
1 parent 7fb4763 commit 9719695

2 files changed

Lines changed: 716 additions & 0 deletions

File tree

‎docs/MAMBA3-PATHC-VS-CPPMEGA.md‎

Lines changed: 108 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,108 @@
1+
# Mamba3 / M2RNN — path_c (ours) vs cppmega per-region compare (gb10 sm_121)
2+
3+
**Status: SKELETON — numbers pending the GB10 phase run.** This doc records the
4+
methodology and the table shape; the GB10 phase fills the MEASURED columns by
5+
running `scratch/mamba3_m2rnn_compare.py` (single-owner, serial). RULE #1: every
6+
cell is MEASURED on gb10 or explicitly `ABSENT`/`NOT-RUNNABLE` with a reason — no
7+
extrapolated or fabricated speedup.
8+
9+
## Why this exists (lever 5)
10+
11+
The SAME mamba3 + m2rnn model is faster in **cppmega** than in **cppmega.mlx
12+
path_c**:
13+
14+
- path_c §17 (`docs/RELAX-GRAPH-VS-MEGATRON.md`): GO **≈907 tok/s @8L / ≈298
15+
@28L** (Relax-VM graph: real gridded CUDA SSD fwd + re-gridded gridded bwd,
16+
carrying an abstract numpy adam/loss as the dominant remaining term).
17+
- the SAME model in cppmega on gb10: **~3692 tok/s @ ~4437 ms/iter** (tensorwise),
18+
torch + CUDA-graph, fully device-resident, where mamba+m2rnn are only ~24% of
19+
the iter (nam56r nsys: MIMO 17.5% + M2RNN 6.2%) and the step is launch-amortized.
20+
21+
The prior READ-ONLY investigation ranked the cause as a **runtime/execution-engine
22+
gap, NOT the SSD scan**. This bench measures that claim with per-region numbers.
23+
24+
## What the bench times (SAME problem config, apples-to-apples)
25+
26+
| arm | region | what runs | source |
27+
| --- | --- | --- | --- |
28+
| OURS (path_c) | F0 | `mamba3_chunk_precompute` (grid, no scan dep) | `mamba3_chunked_precompute_core.py` |
29+
| OURS | F1 | `mamba3_inter_chunk_recur` (O(S/C) recurrence) | `mamba3_chunked_precompute_core.py` |
30+
| OURS | F2 | `mamba3_chunk_scan_combine` (grid scan+combine → y) | `mamba3_chunked_scan_core.py` |
31+
| OURS | B2 | `mamba3_chunk_scan_combine_bwd` (grid) | `mamba3_chunked_backward_core.py` |
32+
| OURS | B1 | `mamba3_inter_chunk_recur_bwd` (O(S/C) reverse) | `mamba3_chunked_backward_core.py` |
33+
| OURS | B0 | `mamba3_chunk_precompute_bwd` (grid) | `mamba3_chunked_backward_core.py` |
34+
| OURS | fwd_chain / bwd_chain | F0→F1→F2 / B2→B1→B0 back-to-back, host-staged | — (the launch-overhead term) |
35+
| cppmega | mamba fwd/bwd | `mamba_chunk_scan_combined` (Megatron SSD Triton — SAME SSD math) | `mamba_ssm.ops.triton.ssd_combined` |
36+
| cppmega | m2rnn fwd/bwd | `m2rnn_scan_triton` (fused Triton M2RNN) | `cppmega/megatron/m2rnn_triton.py` |
37+
| OURS | m2rnn | **NOT-RUNNABLE-ON-CUDA** — Path-C M2RNN is Metal/MSL-only, no CUDA twin | `cppmega_mlx/nn/_tilelang/m2rnn_path_c.py` |
38+
39+
**Honest gap (RULE #1):** our M2RNN has NO CUDA implementation; on gb10 there is
40+
no ours-m2rnn kernel to time. The bench reports this as `NOT-RUNNABLE-ON-CUDA`
41+
and times the cppmega m2rnn alone (absolute region cost), so the §17 finding —
42+
that mamba+m2rnn are a MINORITY of the iter — can still be confirmed.
43+
44+
## The two questions the numbers answer
45+
46+
1. **Is the loser the KERNEL or the HOST/LAUNCH STAGING?**
47+
- `ours_fwd_host_stage_ms = fwd_chain − (F0+F1+F2)`; if this is a large
48+
fraction of `fwd_chain`, the gap is launch/host staging (H1/H4), not the scan.
49+
- `fwd_kernels_only_vs_fused = (F0+F1+F2) / cppmega_fused_fwd`; if ≈1, our
50+
kernels are competitive and the gap is elsewhere; if ≫1, the kernel is the loser.
51+
2. **How big is the m2rnn region vs the mamba region?** Confirms the §17 claim
52+
that the mamba/m2rnn ops are a minority of the iter (so the engine, not the
53+
op, is the lever).
54+
55+
## MEASURED results (gb10) — TO BE FILLED BY THE GB10 PHASE
56+
57+
Config: `--prod` = local_gb10_quarter mamba tile S=4096 c=64 g=8 H=112 P=64 N=64
58+
(== §17). m2rnn shape S=4096 H=8 K=128 V=128. bs1 first; `--bs4` for the 16384-tok
59+
batch axis.
60+
61+
### Per-region median ms/call
62+
63+
| region | OURS ms | cppmega ms | ours/cppmega |
64+
| --- | --- | --- | --- |
65+
| F0 | _pending_ | — | — |
66+
| F1 | _pending_ | — | — |
67+
| F2 | _pending_ | mamba fwd _pending_ | _pending_ |
68+
| fwd_chain (F0→F1→F2) | _pending_ | — | — |
69+
| fwd_host_stage (chain − Σregions) | _pending_ | — | — |
70+
| B2 | _pending_ | — | — |
71+
| B1 | _pending_ | — | — |
72+
| B0 | _pending_ | mamba bwd−fwd _pending_ | _pending_ |
73+
| bwd_chain (B2→B1→B0) | _pending_ | — | — |
74+
| bwd_host_stage | _pending_ | — | — |
75+
| m2rnn fwd | NOT-RUNNABLE | _pending_ | n/a |
76+
| m2rnn bwd | NOT-RUNNABLE | _pending_ | n/a |
77+
78+
### Verdict (filled from `RESULT.attribution.verdict`)
79+
80+
- `fwd_loser`: _pending_ (HOST/LAUNCH-STAGING vs KERNELS)
81+
- `ours_m2rnn`: NOT-RUNNABLE-ON-CUDA (Metal-only Path-C; UNMEASURABLE on gb10)
82+
- mamba vs m2rnn region split: _pending_
83+
84+
## §17 anchors (for cross-check, already MEASURED)
85+
86+
- F2 forward scan = **0.980 ms** vs serial 6.56 s (≈6694×); gridded fwd chain
87+
**7.231 ms/call** (§15, fp16 parity 4.746e-04 PASS).
88+
- gridded backward: B2 re-gridded **2484 → 334.6 ms (7.42×)**, chain **2601 →
89+
447.8 ms (5.81×)**, 1.018× faster than the 456 ms numpy backward; ALL 8 grads
90+
pass the 1e-3 gate (dD 2.48e-5).
91+
- cppmega same-model gb10: **3692 tok/s @ 4437 ms/iter** (tensorwise), 25.7 GB.
92+
93+
## How to run (GB10 phase, single-owner serial)
94+
95+
```
96+
# ensure EXCLUSIVE GPU ownership + >105 GB free; SIGTERM (not -9) above 113 GB
97+
# ensure_nvrtc_builtins_path() BEFORE torch import (cppmega_mlx/_gb10_nvrtc_env.py)
98+
PYTHONPATH=<cppmega_mlx>:<tvm/python>:<tvm-ffi/python> \
99+
TVM_LIBRARY_PATH=<tilelang/build/lib> \
100+
CPPMEGA_ORIG_ROOT=<cppmega original checkout on gb10> \
101+
/home/dave/cppmega-venv/bin/python scratch/mamba3_m2rnn_compare.py --prod
102+
# then --prod --bs4 for the 16384-tok batch axis.
103+
# grep '^RESULT ' <log> | sed 's/^RESULT //' | python -m json.tool # parse
104+
```
105+
106+
If `mamba_ssm` (the Megatron SSD kernel) or `triton` is absent, the cppmega arm
107+
is recorded as a hard error in `RESULT.hard_errors` and printed loudly (RULE #1:
108+
surfaced, not swallowed) — install the dep and re-run; ours-mamba still times.

0 commit comments

Comments
 (0)