feat(moe): layout-API MXFP4 (a4w4/a8w4) MoE gemm

.humanize/kernel-agent/gemm2-att-analysis.md

@@ -0,0 +1,168 @@
+# gemm2 ATT stall attribution: ours-best vs aiter, and the B-prefetch-hoist fix
+
+rocprofv3 **Advanced Thread Trace (ATT)** instruction-timeline analysis of the
+isolated GPT-OSS MoE gemm2 (down-proj) kernel, decomposing where OUR gemm2 loses
+time vs aiter's cktile a16w4 stage2, and the fix that closes part of the gap.
+
+- GPU gfx950 (MI350X), `HIP_VISIBLE_DEVICES=7`, cold (`FLYDSL_RUNTIME_ENABLE_CACHE=0`).
+- Shape: GPT-OSS M=128, model=inter=2880->pad 3072, E=128, topk=4, a4w4 (primary)
+  / a8w4, mxfp4 w2, per_1x32 E8M0 scale. Identical fixed-seed inputs + routing on
+  both sides (`same_input_parity.build_shared_inputs`), per-expert histogram
+  identical -> identical 554.5 MB weight+scale traffic (established, not re-derived).
+- Tool: **rocprofv3 1.1.0 ATT** (`advanced_thread_trace: true`, `att_target_cu: 1`,
+  full SE/SIMD mask), `FLYDSL_DEBUG_ENABLE_DEBUG_INFO=1` for source mapping (ours
+  99% src-mapped; aiter's prebuilt CK binary has no debug info -> ASM-level only).
+  Analyzer: `.claude/skills/kernel-trace-analysis/scripts/hotspot_analyzer.py`.
+
+Configs traced:
+- (i)  ours BEST = BM32 + reduce + G (`MXFP4_G2_KSTAGES=2` B-prefetch, NT w2).
+       kernel `gemm2_a4w4_port_h3072_imax8192_bm32_nt_reduce_tk4_pad_g2ks2_v2`.
+- (ii) ours BM16 + atomic (aiter-tile-aligned). kernel `..._bm16_nt_atomic_pad_v2`.
+- (iii) aiter BM16 flatmm (`.../aiter_gemm2_isolated.py`), `ck_tile::MoeFlatmmKernel`.
+
+Isolated device-kernel times (rocprofv3 kernel-trace, median over ~50 in-run
+dispatches, this GPU/env):
+
+| config | gemm2 median us | eff BW | ratio vs aiter |
+|---|---:|---:|---:|
+| aiter BM16 flatmm | **93.0** | 5.96 TB/s | 1.00x |
+| ours BEST (BM32+reduce+G) | 97.8 | 5.67 TB/s | 1.055x |
+| ours BM16 + atomic | **169.5** | 3.27 TB/s | 1.82x (dead end) |
+
+(The ~1.055x here is tighter than the doc's 1.10-1.13x because ours-BEST already
+carries G's B-prefetch and this GPU runs slightly hotter; the *diagnosis* is the
+same.)
+
+---
+
+## 1. Stall breakdown (% of total stall cycles, steady-state dispatch, single CU)
+
+| stall class | ours BEST | aiter | ours BM16-atomic |
+|---|---:|---:|---:|
+| **VMEM-wait** (`s_waitcnt vmcnt`) | 52.6% | 57.7% | 43.5% |
+| **VMEM-load** (buffer_load issue/back-pressure) | 30.4% | 21.8% | 45.2% |
+| **barrier** (`s_barrier`) | **11.3%** | 6.6% | 7.8% |
+| MFMA | 1.8% | 0.5% | 0.1% |
+| LDS (ds_read/ds_write) | 0.4% | 1.3% | 0.1% |
+| LDS/SMEM-wait (`lgkmcnt`) | 1.5% | 1.8% | 0.6% |
+| other | 2.1% | 10.2% | 2.6% |
+| **total-stall / total-cycles** | **94.1%** | **84.3%** | ~96% |
+
+Occupancy on the traced (busy) CU: **both ours and aiter run 4 waves/SIMD,
+VGPR-bound** (ours VGPR 128, aiter 104; both hit 4 waves; avg waves-in-flight ~3.9
+each). So at the instruction level on a hot CU the LDS-blocks/CU story is NOT the
+binding limiter — both are VGPR-capped at 4 waves and **VMEM(weight)-bound**.
+
+### The dominant differences
+
+1. **Both kernels are weight-VMEM-bound** — ours 83% (VMEM-wait+load), aiter
+   79.5%. HBM is near-saturated; there is no LDS-wait, atomic-store, or MFMA
+   bottleneck. The reduce epilog contributes <2% (LDS-wait + a handful of store
+   waitcnts) — **not epilog-bound, not LDS-bound, not atomic-bound.**
+
+2. **Ours-BEST's distinguishing excess is barrier idle: 11.3% vs aiter's 6.6%**
+   (ours 3 barriers costing 1.18M stall cycles, ~394K each; aiter 6 barriers
+   costing 631K, ~105K each). Ours' single per-K-iteration `gpu.barrier()`
+   (guarding the A-LDS triple-buffer ring, `mxmoe_gemm_v2.py:1385`) is a full-block
+   sync that **absorbs the per-wave VMEM-latency imbalance**: waves whose weight
+   loads returned late drag all 4 waves at the barrier. aiter processes 2 K-blocks
+   per iteration with finer 16x16x32 MFMAs, so its barriers are both less frequent
+   relative to compute and cheaper per hit.
+
+3. **BM16-atomic confirms the pipeline-underfeed root cause** (supporting
+   evidence, not the target). Halving BM to 16 halves our wide-MFMA density (8 vs
+   16 16x16x128 MFMAs / K-tile) -> **VMEM-load back-pressure explodes to 45.2%**
+   (queue-full: the weight loads can't be drained because there isn't enough
+   compute to hide them) and the kernel balloons to 169us. This is why BM16 is a
+   dtype dead-end for our wide scaled-fp4 MFMA: fewer/bigger MFMAs under-feed the
+   weight stream.
+
+**Dominant class = pipeline / VMEM-wait (weight stream not back in time), with
+the ours-specific residual concentrated in barrier idle coupled to that VMEM
+imbalance. NOT LDS, NOT atomic, NOT epilog.**
+
+---
+
+## 2. The fix: hoist the B-weight prefetch above the mainloop barrier (`g2_bhoist`)
+
+The G 2-stage B pipeline (`g2_kstages==2`) already prefetches the next K-tile's B
+weight+scale one tile ahead, but it issued that prefetch **after** the
+per-iteration `gpu.barrier()`. B is a GMEM->register stream with **no LDS
+dependency** — the barrier only orders the A-LDS ring, so it does not need to gate
+the weight loads. Moving the prefetch **above** the barrier puts the long-latency
+weight loads into the VMEM queue *before* the block synchronizes, so the weight
+fetch overlaps the barrier wait instead of starting after it.
+
+- Opt-in `g2_bhoist` param on `gemm2_body_v2` (default `False` = byte-identical),
+  env `MXFP4_G2_BHOIST=1`, threaded through the dispatcher cache key + kernel-name
+  tag (`_bhoist`). No-op unless `g2_kstages==2`. gemm1 untouched.
+- The prefetch body was factored into a `prefetch_next_b(kt_rt)` closure called
+  either before (hoist) or after (default) the barrier via `const_expr`.
+
+### Effect (ATT, same dispatch index, ours BEST vs BEST+bhoist)
+
+| stall class | BEST | BEST + bhoist |
+|---|---:|---:|
+| VMEM-wait | 52.6% | **34.9%** |
+| VMEM-load | 30.4% | 46.8% |
+| barrier | 11.3% | 12.4% |
+| total stall (M cycles) | 10.46M | **10.17M** |
+
+The hoist does exactly what the diagnosis predicted: **VMEM-wait drops
+52.6% -> 34.9%** (the `s_waitcnt vmcnt` on the weights waits less because the
+loads were issued earlier), shifting into VMEM-load queue-occupancy (loads now
+in flight during the barrier). Total stall cycles drop; the profile moves toward
+aiter's "loads always flowing" shape.
+
+### Measured gemm2 (median device-kernel us, ~50 dispatches, cold)
+
+| config | a4w4 median us | eff BW | ratio vs aiter (93.0us) |
+|---|---:|---:|---:|
+| ours BEST | 97.8 | 5.67 TB/s | 1.055x |
+| **ours BEST + bhoist** | **95.6** | **5.80 TB/s** | **1.028x** |
+
+a4w4: **97.8 -> 95.6 us (-2.2%)**, min 96.2 -> 93.5 us; ratio **1.055x -> 1.028x**.
+
+a8w4: **100.5 -> 100.8 us (neutral, within noise)** — a8w4's heavier A path
+(`KH_TILE_A=256`, 2x A-LDS DMA) is not as purely weight-VMEM-wait-bound, so the
+hoist has nothing to hide there. Neutral is acceptable since it is opt-in.
+
+---
+
+## 3. Correctness + byte-identical + gemm1-untouched proofs
+
+- **Correctness (cold, real 2880 dims):** a4w4 cos = **0.9910** (>0.85),
+  a8w4 cos = **0.9996** (>0.95) — bitwise the same cos as without bhoist (the
+  reorder does not change the math, only issue order).
+- **Byte-identical default (AC-3):** the shipped default (`MXFP4_G2_BHOIST` unset
+  -> `g2_bhoist=False`, and the shipped `g2_kstages=1` default) emits no `_bhoist`
+  tag and `prefetch_next_b` stays in its original position. ISA md5 of the default
+  `g2ks2` kernel is **identical** across the edited tree vs the pre-change source:
+  `18b5e4d7ed041e45e2705f333a3e2f53` (both). (g2_kstages=1, the true shipped
+  default, is untouched code.)
+- **gemm1 untouched:** `git diff` touches only the gemm2 body's `g2_kstages==2`
+  loop and the gemm2 dispatcher path; no gemm1 line changed.
+
+---
+
+## 4. Verdict on the residual
+
+After bhoist, ours-BEST is **1.028x** aiter (95.6 vs 93.0 us, 5.80 vs 5.96 TB/s).
+The ATT shows the kernel is now **VMEM-load-queue-bound (46.8%)** — i.e. HBM is
+essentially saturated, loads are always in flight, and the remaining ~2.8% gap is
+the last sliver of aiter's finer 2-K-block/16x16x32 interleave keeping the queue
+marginally fuller. That residual is small and hard to close without adopting
+aiter's bf16-A / finer-MFMA structure (a dtype regression for us, per the
+deepdiff) or a deeper (kstages 3) B pipe that would add VGPR pressure and risk
+dropping below 4 waves/SIMD. The dominant, actionable stall (weight-VMEM-wait +
+barrier-coupled imbalance) has been attacked; further gains are in diminishing-
+returns territory against a near-saturated HBM.
+
+## Artifacts
+
+- ATT dirs: `/tmp/att_ours_best/`, `/tmp/att_bhoist/`, `/tmp/att_bm16/`,
+  `/tmp/att_aiter/` (ui_output_agent_*_dispatch_* with code.json + per-wave traces).
+- Code: `kernels/mxmoe_gemm_v2.py` (`g2_bhoist` param + `prefetch_next_b` closure),
+  `kernels/moe_dispatcher.py` (env `MXFP4_G2_BHOIST`, cache key, `_bhoist` tag).
+- Drivers: `.humanize/kernel-agent/same_input_parity.py` (ours),
+  `.humanize/kernel-agent/aiter_gemm2_isolated.py` (aiter).

.humanize/kernel-agent/gemm2-spatial-partitioner.md

@@ -0,0 +1,192 @@
+# gemm2 channel balance via aiter's portable spatial tile partitioner (replaces xcd)
+
+Port aiter's ACTUAL, XCD-count-INDEPENDENT gemm2 block->tile mapping — the
+`GemmSpatiallyLocalTilePartitioner` grouped 2D rasterization — as our gemm2
+block->(m_block_idx, n_block_idx) map, and DROP the non-portable explicit-8-XCD
+swizzle (`g2_xcd`). This is the true aiter alignment for HBM channel balance.
+
+Setup: gfx950 MI350X, `HIP_VISIBLE_DEVICES=3`, cold
+(`FLYDSL_RUNTIME_ENABLE_CACHE=0`), GPT-OSS M=128, model=inter=2880->pad 3072,
+E=128, topk=4, mxfp4 w2, per_1x32 E8M0 scale. Ours-best base config (worktree
+branch `rlcr/g2-spart` @ 947c766c) = BM32 + reduce + g2ks2 (B-prefetch) + bhoist
++ apf. Drivers: `same_input_parity.py` (ours), `aiter_gemm2_isolated.py` (aiter).
+
+Prereq reading: `gemm2-mem-pattern-align.md` (the measured channel-imbalance
+root cause + the xcd result to beat), `cktile-gemm2-deepdiff.md`.
+
+## 0. THE KEY FINDING: aiter's partitioner is GroupNum=1, M01=1 — a no-op
+
+aiter's shipped gemm2 instantiates
+`GemmSpatiallyLocalTilePartitioner<CodegenFlatmmShape, GroupNum, M01>` with
+**`GroupNum=1`, `M01=1`** — hard-coded in `FlatmmConfig` for EVERY instance, no
+override anywhere:
+
+```
+csrc/ck_tile_gemm_moe_2stages/include/moe_cktile2stages_common.cuh:58-59
+    static constexpr int TileParitionerGroupNum = 1;
+    static constexpr int TileParitionerM01      = 1;
+    ...:102-105  using TilePartitioner = GemmSpatiallyLocalTilePartitioner<shape, 1, 1>;
+```
+
+`MoeFlatmmKernel::operator()` calls `TilePartitioner{M,N}.GetOutputTileIndex(blockIdx.x)`
+(`moe_flatmm_kernel.hpp:757-762`). I traced `GetOutputTileIndex`
+(`gemm_tile_partitioner.hpp:274-360`) at `GroupNum=1, M01=1` by hand and by
+numeric replay over realistic grids (`/tmp/spart_check.py`):
+
+- `group_size = ceil(M0*N0/1) = M0*N0`, `big_group_num = 1`, so
+  `remap_block_1d_id == block_1d_id` (the big-group interleave is identity).
+- `M0_mod_M01 = 0`, `M01_adapt = 1`, `idx_M01 = 0`, so
+  `return (idx_M0, idx_N0) = (block_1d_id / N0, block_1d_id % N0)`.
+
+**At (1,1) the partitioner is byte-identical to the naive m-major linear map we
+already use** (`m_block_idx = bx // num_n_blocks`, `n_block_idx = bx % num_n_blocks`).
+So aiter's spatial partitioner is NOT the source of its channel balance for this
+instance — porting it verbatim would be a no-op (still CV 9.7%). Verified:
+`get_output_tile(b,M0,N0,1,1) == (b//N0, b%N0)` for all b over M0 in
+{37,40,125,512}, N0=12.
+
+Where aiter's balance actually comes from: its gemm2 tiles N at 128 (vs our 256),
+runs a persistent grid (Kind2 grid 983040), and reads bf16 A — a different tiling
+regime, not the partitioner grouping. The productive port is therefore the
+GENERAL `GetOutputTileIndex` parameterized by (GroupNum, M01), sweeping
+NON-TRIVIAL values that actually rasterize spatially-locally and rebalance
+channels — which is exactly what the partitioner is designed to do, just with the
+grouping ENABLED (GroupNum/M01 > 1) rather than aiter's disabled (1,1).
+
+## 1. The port (`g2_spart`, opt-in, default byte-identical)
+
+Opt-in `g2_spart` param on `compile_gemm2_a4w4_port`, env `MXFP4_G2_SPART`,
+encoded `GroupNum*100 + M01` (e.g. `402` = GroupNum4, M01=2); `0`/unset = off =
+byte-identical naive linear grid. Threaded through the gemm2 dispatcher cache key
++ kernel-name tag (`_spart{G}x{M01}`). gemm1 untouched (the remap lives entirely
+in `moe_dispatcher.py`'s gemm2 one-shot path; `mxmoe_gemm_v2.py` unchanged).
+
+`_spart_output_tile_index(block_1d_id, M0, N0, GroupNum, M01)` is a faithful DSL
+port of `GetOutputTileIndex`'s else-branch (M0=total_m_blocks runtime;
+N0=num_n_blocks, GroupNum, M01 compile-time): group_size / big_group_num remap,
+then the M01 spatially-local M-window re-tile. It emits the same grouped
+rasterization; consecutive block ids stay spatially local in the M0xN0 grid so
+concurrent blocks' weight fetches spread across HBM channels — WITHOUT any XCD
+awareness (no hard-coded 8; portable across GPUs).
+
+Because the remap needs `M0 = total_m_blocks` (runtime, from the cumsum), the
+spart path reads the cumsum FIRST (losing the default's A-prologue/cumsum overlap)
+then feeds `unit_bx = m_block_idx*num_n_blocks + n_block_idx` to the A prologue +
+`run_unit` — same tradeoff the xcd path made. `gemm2_body_v2` re-decodes `unit_bx`
+linearly back to the same (m,n), so the port is consistent with the body.
+
+**Bijection over [0, M0*N0):** verified for all candidate (GroupNum,M01) over
+M0 in {1,2,3,37,40,125,512}, N0=12 — every (m,n) tile computed exactly once, no
+dropped/duplicated tiles (`/tmp/spart_check.py`).
+
+## 2. MEASURED channel distribution (rocprofv3 TCC_EA0_RDREQ, 128 instances)
+
+Per-channel HBM read requests over the 128 TCC instances (16 channels x 8 XCC),
+median over the replayed a4w4 gemm2 dispatches (warmup dropped), same GPU3/session.
+Method identical to `gemm2-mem-pattern-align.md` §4b.
+
+| config | total EA reads | per-chan mean | min | max | **CV** | **max/min** |
+|---|---:|---:|---:|---:|---:|---:|
+| naive (linear, spart off) | 4,857,351 | 37,948 | 31,495 | 40,198 | **9.68%** | **1.28x** |
+| spart 202 (G2,M01=2) | 4,557,421 | 35,605 | 33,442 | 37,764 | 5.75% | 1.13x |
+| **spart 402 (G4,M01=2)** | **4,501,146** | 35,165 | 34,992 | 35,364 | **0.32%** | **1.01x** |
+| spart 404 (G4,M01=4) | 4,509,433 | 35,230 | 35,070 | 35,497 | 0.35% | 1.01x |
+| spart 802 (G8,M01=2) | 4,508,507 | 35,223 | 34,897 | 35,483 | 0.48% | 1.02x |
+| spart 804 (G8,M01=4) | 4,529,698 | 35,388 | 35,050 | 35,803 | 0.55% | 1.02x |
+| spart 1204 (G12,M01=4) | 4,563,093 | 35,649 | 35,430 | 35,977 | 0.44% | 1.02x |
+| **xcd4 (prior, doc §5)** | 4,493,066 | — | — | — | **0.64%** | **1.02x** |
+| **AITER (same session)** | **4,436,910** | 34,663 | 34,590 | 34,759 | **0.18%** | **1.00x** |
+
+**spart 402 (GroupNum=4, M01=2) is the winner: CV 9.68% -> 0.32%, max/min 1.01x,
+total reads 4.86M -> 4.50M.** It BEATS the prior xcd4 path on channel balance
+(0.32% vs 0.64% CV) and approaches aiter's 0.18%. Total HBM reads land within
+1.4% of aiter's 4.44M (was 9.5% over) — the imbalance was forcing the extra
+fetches; even distribution removes them, same as the xcd path. The portable
+partitioner MATCHES-OR-BEATS the XCD swizzle as the channel-balance mechanism.
+
+## 3. MEASURED perf (gemm2 isolated, median-of-5 cold, same GPU3 session)
+
+Interleaved 5 rounds of naive / spart402 / aiter to control gfx950 clock drift.
+Absolute us is ~12% higher than `gemm2-mem-pattern-align.md`'s numbers because
+this GPU/session runs hotter (see +-14% gfx950 clock noise in MEMORY); the
+RELATIVE deltas and the aiter ratio measured IN THIS SESSION are load-bearing.
+
+| config | a4w4 median us | a8w4 median us | ratio vs aiter (95.3) |
+|---|---:|---:|---:|
+| aiter cktile a16w4 | 95.3 | — | 1.00x |
+| naive (linear) | 109.4 | 112.0 | 1.148x |
+| **spart 402** | **108.3** | **109.9** | **1.136x** |
+
+- a4w4: **109.4 -> 108.3 us (-1.0%)**, ratio vs aiter 1.148x -> 1.136x.
+- a8w4: **112.0 -> 109.9 us (-1.9%)** (its heavier A path is more
+  channel-sensitive, same pattern the xcd path showed).
+
+raw us (GPU3, this hotter session):
+```
+naive a4w4:    109.1 109.3 109.4 109.5 109.8 -> med 109.4
+spart402 a4w4: 108.3 107.6 108.3 108.5 108.2 -> med 108.3
+naive a8w4:    112.0 112.0 111.8 112.0 111.3 -> med 112.0
+spart402 a8w4: 110.2 109.7 110.0 109.9 109.4 -> med 109.9
+aiter a4w4:    95.1  95.3  95.5  95.3  95.3  -> med 95.3
+```
+
+### spart vs xcd: matches on mechanism, perf comparison caveat
+
+On the SESSION-INDEPENDENT channel metric spart402 (CV 0.32%, 4.50M reads)
+MATCHES-OR-BEATS xcd4 (CV 0.64%, 4.49M reads) — the two are equivalent on total
+HBM reads and spart is tighter on CV. The doc §5 xcd perf (94.5us, 0.993x vs
+aiter) was measured in a COOLER session (its naive baseline was 97.1us there vs
+109.4us here), so the absolute us are not directly comparable across sessions. A
+same-session xcd re-measurement was attempted but the ported xcd swizzle block
+(from commit 93ad0bd4, a different base tree) faulted (hipErrorIllegalAddress) on
+this base config and was reverted unused — xcd is the mechanism being DROPPED, so
+it is not carried in this tree. The channel-balance equivalence is established on
+the portable metric; the smaller absolute perf gain here reflects the
+memory-headroom-compressed hotter session, not a weaker mechanism (same +9.1%
+over-fetch is removed, total reads land at aiter parity).
+
+## 4. Correctness + byte-identical default + gemm1-untouched proofs
+
+- **Correctness (cold, real 2880 dims, spart402):** a4w4 cos = **0.9910**
+  (>0.85), a8w4 cos = **0.9996** (>0.95) — identical to baseline (bijective block
+  permutation, same math). Thresholds NOT weakened.
+- **Byte-identical default (AC-3):** with `MXFP4_G2_SPART` unset the emitted
+  default gemm2 kernel (`gemm2_a4w4_port_h3072_imax8192_bm32_reduce_tk4_pad_g2ks2_bhoist_apf_v2`,
+  no `_spart` tag) is **md5-identical** with the spart code present vs stashed:
+  `8263440a146105e9b138b109296a8648` both ways (empty diff). No extra IR emitted
+  when off.
+- **gemm1 untouched:** `git diff` touches ONLY `kernels/moe_dispatcher.py`
+  (+98/-2), all in the gemm2 dispatch path (`_spart_output_tile_index` helper,
+  `compile_gemm2_a4w4_port` param/env/tag, the one-shot `elif` remap branch,
+  `get_g2` cache key). No `gemm1` line changed; `mxmoe_gemm_v2.py` unchanged.
+  Python style gate passes (black + ruff).
+
+## 5. Verdict
+
+aiter's REAL gemm2 tile partitioner is `GemmSpatiallyLocalTilePartitioner` with
+**GroupNum=1, M01=1, which degenerates to the naive linear map** — it is NOT what
+balances aiter's channels (that comes from aiter's N=128 tiling + persistent grid
+regime). Porting the GENERAL partitioner with the grouping ENABLED
+(**GroupNum=4, M01=2 = `MXFP4_G2_SPART=402`**) rebalances OUR channels from CV
+9.68% to **0.32%** (beating the prior xcd4's 0.64%, approaching aiter's 0.18%) and
+cuts total HBM reads to 4.50M (aiter parity 4.44M), with a4w4 -1.0% / a8w4 -1.9%
+gemm2 time in this session, correctness held, default byte-identical, gemm1
+untouched. This portable, XCD-count-INDEPENDENT spatial partitioner REPLACES the
+non-portable `g2_xcd` explicit-8-XCD swizzle as the channel-balance mechanism =
+the true aiter alignment.
+
+## Artifacts
+
+- Code: `kernels/moe_dispatcher.py` (`_spart_output_tile_index` DSL port of
+  `GetOutputTileIndex`; `g2_spart` param + env `MXFP4_G2_SPART` + `_spart{G}x{M01}`
+  tag on `compile_gemm2_a4w4_port`; one-shot `elif` remap branch; `get_g2` cache
+  key). gemm1 and `mxmoe_gemm_v2.py` untouched.
+- aiter source: `moe_cktile2stages_common.cuh:58-59,102-105` (GroupNum=1/M01=1);
+  `ck_tile/ops/gemm/kernel/gemm_tile_partitioner.hpp:274-360`
+  (`GetOutputTileIndex`); `moe_flatmm_kernel.hpp:757-762` (its use).
+- PMC captures: `/tmp/g2prof/spart_{naive,202,402,404,802,804,1204,aiter}/`
+  (rocprofv3 json, TCC_EA0_RDREQ). Parser `/tmp/g2prof/perchan.py`.
+- Timing: `/tmp/g2prof/times.csv` + `/tmp/g2prof/full_time2.log`.
+- Partitioner math check: `/tmp/spart_check.py` (degeneration + bijection).
+- Drivers: `.humanize/kernel-agent/same_input_parity.py`,
+  `aiter_gemm2_isolated.py`.