|
2 | 2 |
|
3 | 3 | All notable UI-surface changes since the V7 honest-closure pass. |
4 | 4 |
|
| 5 | +## V7-D-block: TileLang→Metal kernel epic closure (2026-05-23 14:30 GMT+2) |
| 6 | + |
| 7 | +D-block (TileLang→Metal kernels, 8 subtickets + Path B adapter + epic |
| 8 | ++ breakdown task = 13 tickets) closed as **stale bd vs reality**. |
| 9 | + |
| 10 | +- **All 10 kernels shipped** in `cppmega_mlx/nn/_tilelang/`: |
| 11 | + `topk_selector` (Path B + Path C, 923 LOC), `sparse_mla` / |
| 12 | + `sparse_mla_path_c` (BF16 fwd+bwd), `sparse_mla_fp8` / |
| 13 | + `sparse_mla_fp8_path_c` (FP8 fwd+bwd), `sparse_mla_blockscaled` / |
| 14 | + `sparse_mla_blockscaled_path_c` (e8m0 fwd+bwd), |
| 15 | + `mamba3` / `mamba3_path_c` (fwd+bwd via Path C), |
| 16 | + `_mamba3_helpers_tilelang.compute_dacs_segsum` (Triton-free). |
| 17 | +- **Path B adapter** (`0ce`): `TileLangMetalAdapter` + |
| 18 | + `wrap_tilelang_metal_kernel` in |
| 19 | + `cppmega_mlx/nn/_tilelang/_mlx_runtime.py:628-876`. Takes TileLang |
| 20 | + artifact → extracts MSL → `mx.fast.metal_kernel`. |
| 21 | +- **Test fix**: `tests/test_path_c_sparse_mla_apply_emit_indent.py` |
| 22 | + was calling `_mamba3_fp8_train_acceptance_region(model_config=…)` |
| 23 | + but the signature was simplified to `include_backward` only. |
| 24 | + Dropped the stale kwarg; 3 tests now PASS. |
| 25 | +- **Regression evidence**: 407/407 kernel tests PASS across 12 |
| 26 | + test files at commit `77a1404`. bd open count: 57 → 44. |
| 27 | + |
| 28 | +Closed tickets: `cppmega-mlx-za1`, `za1.1`, `za1.2`, `za1.3`, `za1.4`, |
| 29 | +`za1.5`, `za1.6`, `za1.7`, `za1.8`, `za1.9`, `za1.10`, `za1.11`, |
| 30 | +`0ce`. |
| 31 | + |
5 | 32 | ## V7-M0.7 + M0.3 closure (2026-05-23 06:00 GMT+2) |
6 | 33 |
|
7 | 34 | - **M0.7 (resumable training) ✅ green** — `stage_train` now |
|
0 commit comments