MFU Gap Investigation: "Expected 50%, Getting 25%"

A staff-level model answer: solve the incident, then design the repeatable process. Hardware assumptions, flagged as such throughout: Trillium-class chip ≈ 918 TFLOP/s bf16 dense, 32 GB HBM @ 1.64 TB/s, 256-chip ICI pod; effective per-chip collective bandwidth for the patterns here assumed ~12 GB/s (a measured-once fleet prior, not a spec).

1 · Clarifying Questions and Framing

Question	Why it changes everything	Assumed answer
Training or inference?	Different denominators, different pathologies, different tools	Training (primary track); inference variant noted below
If inference — prefill or decode MFU?	Quoting one MFU for inference is already a red flag: decode is memory-bound by design — its "low MFU" is the roofline, not a bug; MBU is the honest decode metric	n/a here, but I'd say this aloud before accepting the premise
Dense or MoE?	MoE MFU accounting (routed vs dense-equivalent FLOPs) can move the number 10× before any engineering happens	Dense ~70B
Hardware, topology, parallelism?	Sets the ideal-time arithmetic and the collective structure	256 Trillium-class chips, one ICI pod, GSPMD: DP=32 × TP=8, seq 8,192, global batch 256 seqs = 2.10M tokens/step, bf16 + custom fp8 quant ops in FFN
How was "expected 50%" derived?	Folklore vs roofline vs comparable — determines whether the gap is even real	"Comparable model got 50%" — which I will audit in §2, and it partially fails the audit

Inference parallel track (one paragraph, said early to show the map): for a decode workload the investigation reshapes — the metric becomes MBU and goodput-under-SLO; Level 1 collectives are ICI all-to-all/allgather in TP decode; Level 2 becomes KV-cache traffic accounting and batch-occupancy analysis; and "25% MFU" would likely be correct behavior at small batch. Same hierarchy, different instruments; I'll run the training track here.

Arithmetic check — the frame everything hangs on:
FLOPs/step = 6·N·D = 6 × 70e9 × 2.10e6 = 8.81e17. Fleet peak = 256 × 918e12 = 2.35e17 FLOP/s.
Ideal step (100% MFU) = 8.81e17 ÷ 2.35e17 = 3.75 s → 50% MFU = 7.50 s · observed 15.0 s = 25.0% MFU ✓
MFU(step) = 3.75/step_time. Every "MFU point" below converts through this line, so points are non-linear in seconds — the bridge chart in §4 does the conversion honestly.

Staff signal: the two framing moves that are already L6: refusing a single "inference MFU" as a coherent quantity before investigating anything, and demanding the provenance of "expected 50%" — an expectation is a measurement with its own error bars, and auditing it is part of the incident, not an insult to whoever said it.

2 · First 30 Minutes: Interrogate the Metric

Before profiling the system, profile the number. Concretely, in order:

2.1 How is our MFU computed? Numerator: analytic 6ND (model FLOPs from shapes) or traced FLOPs (whatever executed)? It must be analytic — traced FLOPs count rematerialization as useful work and silently absolve recompute overhead (a 30% remat tax vanishes from MFU computed on traced FLOPs). Attention term: with causal masking the true work is ~half the s² rectangle — if the comparable model's MFU used the halved term and ours uses full, that's ~3–4 points of phantom gap on 8K sequences before anything real. Denominator: bf16 dense peak (918), never the sparsity-inflated marketing peak, and never the fp8 peak unless all GEMMs run fp8 (ours are mixed — using fp8 peak would halve reported MFU dishonestly). Wall clock: steady-state step time, excluding init/compile, but including data stalls and amortized checkpointing — because tokens/day is what the fleet buys; a per-step MFU that excludes stalls is a vanity metric. Vocab GEMM and padding FLOPs: in the numerator only if genuinely computed and useful — pad FLOPs are not useful work.

2.2 Triangulate before believing. Three independent lenses must agree: (a) step time 15.0 s vs analytic minimum 3.75 s → 4.0× — large but coherent; (b) tokens/s/chip = 2.10e6/15.0/256 ≈ 547 vs comparable-model ≈ 1,100 → 2.0× worse ✓ consistent with 25 vs 50; (c) HBM BW utilization ≈ 38% average with a periodic near-zero dip (already a clue: compute-starved intervals, §4-iii). If lens (b) had shown parity, the incident would be an accounting bug, and it would be the second time this quarter.

2.3 Audit the expectation. Geometry comparability check against the 50%-model:

Arithmetic check — the expectation is ~3 points wrong on geometry alone:
Comparable: d_model 8,192, head_dim 128, d_ff 28,672 (=112×256 ✓), vocab 131,072 (=512×256 ✓). Ours: identical except head_dim 96 (a GPU-heritage choice) and vocab 151,936 (=593.5×256 ✗).
SDPA GEMMs (QK^T, AV) have inner/lane dim 96 → padded to 128 on the MXU tile: 25% of lanes compute zeros on the ~11% of FLOPs that are SDPA → ~2.7 MFU-point lower ceiling. Vocab GEMM pad waste on ~5% of FLOPs → ~0.3 pts.
Corrected expectation ≈ 47%, not 50. The gap to explain is 25→47; the last ~3 "missing" points were never available without a model-side geometry change. This is the §4 residual, resolved honestly at the start rather than chased at the end.

Staff signal: interrogating the metric before the system — the remat/causal/MoE/padding accounting traps, the three-lens triangulation, and the geometry audit that re-derives the expectation — routinely closes 20–40% of "MFU gaps" without touching the workload. L5 opens the profiler; L6 first establishes that the number and the target are both real.

3 · The Top-Down Decomposition

L0 — time accounting (tool: xprof/Perfetto pod-level timeline; artifact: bucketed step decomposition). Incident numbers: MXU-busy 60% (9.0 s — but §L2 will show a third of it is waste), exposed ICI 24% (3.6 s), infeed/host gaps 10% (1.5 s amortized — visible as periodic, not per-step), framework/gaps 6% (0.9 s). Buckets sum to 15.0 s ✓ — L0 exits.

L1 — communication (artifact: per-collective matrix — name, bytes, count/step, wire-time, exposed-time). Findings pattern: (i) exposed ≫ wire-time → scheduling/serialization, not bandwidth; (ii) a collective that shouldn't exist → accidental resharding from a sharding-annotation mismatch (grep the HLO for all-gather ops whose operand sharding differs only by a transposed mesh axis); (iii) DCN hops appearing in what should be an ICI-only job → placement bug, 10× cost. Incident finding: grad reduce-scatter/all-gather chain wire-time ≈ 2.9 s should hide under 7+ s of backward but runs 3.6 s exposed — every collective queued after a CustomCall the fp8 quant library marks as a scheduling barrier. That's a one-line fix pattern with a 3.6 s payoff.

L2 — compute efficiency (artifacts: per-op top-k table with achieved TFLOP/s and AI; post-optimization HLO fusion audit). Findings pattern: GEMMs at 25–40% of peak → shape/padding (check dims mod 256); duplicate forward kernels inside backward → rematerialization policy surprises; elementwise chains unfused across a custom-op boundary → HBM round-trips. Incident findings: SDPA GEMMs at 34% (head_dim 96 padding — the §2 expectation issue, model-side, parked) and the FFN forward kernels appearing twice in backward: the quant CustomCall is opaque to the remat policy, so XLA recomputes the entire FFN block instead of the cheap ops — +29% executed FLOPs ≈ 2.0 s at achieved rate.

L3 — host & input (artifact: step-time strip chart + host CPU/infeed profile). Incident finding: every 20th step stalls ~21 s → amortized 1.04 s/step: synchronous checkpoint serialization (12 s) + input-shard refill without prefetch (9 s). Periodicity in the strip chart is the tell — mean step time hides it, p95 screams it.

L4 — fleet/hardware (artifact: per-chip step-time distribution, per-link ICI counters). Incident finding: clean — per-chip spread 1.8% (TPU determinism doing its job). On the GPU fleet this level is louder: DVFS spread 5–8% under thermal load, one throttling straggler gates every synchronous step.

Staff signal: the hierarchy's power is its decision rules — descend by bucket size, per-level exit criteria, parking list for sub-2-point items — which make it teachable and repeatable by an L4, which is precisely the point of part (b) of the question. Naming the "collective that shouldn't exist" and "forward kernels twice in backward" patterns shows the artifacts have actually been read before.

Interviewer follow-ups (§3)
Q: How does this hierarchy change on the GPU fleet?
A: Structure identical, instruments and priors shift: L0 runs on nsys/Kineto instead of xprof; L1 swaps ICI for NCCL over NVLink/IB — add NCCL_DEBUG and per-ring analysis, and watch for rail-misrouted traffic; L2 adds kernel-launch overhead and CUDA-graph coverage as first-class buckets (µs launches × thousands matter at small per-GPU work), and cuBLAS heuristic misses replace XLA fusion audits; L4 gets much louder — lock clocks before trusting any number, since DVFS alone can manufacture a 5% "regression." TPU's determinism means n=3 repeats suffice where GPU wants n=10.
Q: What if the profiler perturbs the workload?
A: Measure the observer: step time with and without profiling — xprof tracing overhead is typically <2% but host-side instrumentation can serialize the input pipeline and manufacture the very stall you're chasing. Rules: capture short windows (20–50 steps) not whole runs; confirm any profiler-discovered stall reproduces in unprofiled step-time telemetry (the strip chart comes from lightweight always-on timers, not traces); and never enable per-op host callbacks in the measurement run. If a pathology appears only under profiling, treat the profiler as the suspect.
Q: The gap only appears at 4× scale (1,024 chips). Your 256-chip investigation finds nothing. Now what?
A: The differential is the tool: run the hierarchy at both scales and diff the L0 buckets — scale-dependent gaps live almost exclusively in L1 (collective latency terms that grow with participant count, DCN crossings appearing past one pod, load-imbalance amplification: a 2% straggler is invisible at 256 and gates every step at 1,024) or L3 (input service shared across more hosts). Reproduce at the smallest failing scale to iterate cheaply — bisect the scale axis (256→512→768) like a commit range, and instrument per-participant collective arrival times to catch the straggler class directly.

4 · The Worked Incident — Three Stacked Causes

Order of fixes: cheapest-confirmation-first, because causes (i) and (ii) both involve the quant library and might interact — a flag flip that confirms (i) in an hour de-risks the attribution of everything downstream. Then (ii), then (iii) — which is independent and schedulable in parallel by the data-infra owner.

#	Cause	Found at	Fix	Step time	MFU
—	baseline	—	—	15.00 s	25.0%
i	Collectives serialized behind quant-library CustomCall barrier → 3.6 s exposed ICI	L1 (collective matrix: exposed ≫ wire)	Library flag removes barrier semantics + scheduler reorder; collectives overlap backward	11.36 s	33.0% (+8.0)
ii	Quant op opaque to remat policy → full FFN recompute in backward, +29% executed FLOPs	L2 (fwd kernels ×2 in bwd; HLO remat cluster)	Register op as remat-transparent (pure, cheap-to-save); policy saves quant scales, recomputes only elementwise	9.37 s	40.1% (+7.1)
iii	Every-20th-step stall: sync checkpoint (12 s) + shard refill (9 s) ≈ 1.04 s/step amortized	L3 (strip chart periodicity)	Async multi-tier checkpointing + double-buffered shard prefetch	8.33 s	45.0% (+4.9)
—	Residual: head_dim-96 padding (~2.7 pts, model-side, next revision) + framework floor (~2.3 pts)	§2 audit + L0 remainder	Accepted; expectation corrected to ~47%	—	45.0 vs 47 corrected target

Arithmetic check — the bridge reconciles:
(i) 15.00 − 3.64 = 11.36 s → 3.75/11.36 = 33.0% ✓ (exposed 3.6 s → 0.4 s residual overlap tax kept in framework bucket)
(ii) executed-FLOPs inflation 1.29× on the compute bucket: compute 6.94 s → 6.94/1.29 ≈ 5.38 s? No — recompute removal deletes 2.0 s of the MXU-busy bucket directly (measured as the duplicated-kernel time): 11.36 − 1.99 = 9.37 s → 40.1% ✓ ; cross-check: 2.0 s ≈ 29% of the 6.9 s useful-compute time ✓
(iii) stall amortization: (12 + 9) s ÷ 20 steps = 1.04 s/step: 9.37 − 1.04 = 8.33 s → 3.75/8.33 = 45.0% ✓
Residual to corrected 47% target: 8.33 vs 47% ⇒ 3.75/0.47 = 7.98 s — remaining 0.35 s ≈ framework floor + imperfect overlap; the further 2.7 geometry points need head_dim 128 in the next model revision.

Staff signal: compound causes expected and found (a 2× gap is almost never one villain); fix order justified by interaction risk, not habit; the bridge chart honestly non-linear (points-per-second changes as the baseline improves — most candidates get this wrong); and the residual reconciled against a corrected expectation instead of hand-waved. The measured-recovery column is what makes this an engineering narrative rather than a theory.

Interviewer follow-ups (§4)
Q: Your fix (i) is a flag in someone else's library. They say the barrier is load-bearing for correctness. Now what?
A: Then the 3.6 s is a design cost, not a bug — and the conversation moves to why: if the barrier guards a host-side scale readback, the fix is keeping scales on-device (which fix (ii) partially does); if it guards genuine ordering, we ask for a scoped fence (barrier only against the quant stream, not all collectives). Meanwhile the scheduler can hoist independent collectives above the CustomCall — a compiler-side workaround worth a bug with the trace attached. The staff move is converting "they said no" into a precise statement of what invariant the barrier protects, because that statement usually reveals a narrower fence.
Q: Fix (ii) claims +29% executed FLOPs from remat. Couldn't that just be your remat policy being optimal and memory forcing it?
A: Checkable, and I'd check it: HBM headroom during backward was 9 GB free — remat wasn't memory-forced, it was policy-blindness (the CustomCall's outputs weren't in the saveable set, so the partitioner recomputed the whole producer chain). If memory had been tight, the answer changes to selective remat (save the quant outputs — small — recompute the cheap elementwise) or offloading scales to host; "recompute everything" is almost never the optimal point on the memory-compute frontier, and the frontier is computable from tensor sizes.
Q: Why did comparable-model history not catch the head_dim-96 issue before training started?
A: Because nothing forced the comparison — which is the process gap §6 closes: the pre-launch geometry checklist diffs every new model's dims against MXU/tile alignment and flags mod-256 violations with their estimated MFU cost before a single chip-hour is spent. This incident's ~2.7-point geometry tax was knowable from the config file alone; the review turns tribal knowledge into a gate.

5 · Artifacts

All artifacts share the incident config. Per-layer backward window ≈ (times to scale; 80 layers dominate the step).

Artifact 1 — timeline BEFORE (one backward layer window, ~142 ms): exposed collective + recompute

Artifact 2 — timeline AFTER (same window, ~86 ms): overlapped, no recompute

Artifact 3 — per-op top-10 (before; % of 15.0 s step)

op                                   time%   ach.TF/s   AI(F/B)   roofline verdict
fusion.3120 ffn bwd GEMM (d_ff)      21.4%     498       ~1,800    compute-bound, 54% peak — layout OK, remat pressure
fusion.3101 ffn fwd GEMM             12.9%     501       ~1,800    compute-bound, 55% peak
RECOMPUTE fusion.3101' (dup in bwd)  13.3%     499       ~1,800    ← should not exist: remat-forced duplicate
reduce-scatter.grad (EXPOSED)        16.1%      —         —        wire 2.9 s vs exposed 3.6 s → scheduling, not BW
all-gather.param                      7.9%      —         —        partially exposed, same barrier
sdpa.qk fusion.2200                   5.9%     312       ~1,100    compute-bound, 34% peak ← head_dim 96 padding
sdpa.av fusion.2210                   5.6%     318       ~1,100    same
vocab.gemm fusion.4001                3.1%     471        ~890     54% peak; pad waste ~0.3 MFU pt
attn.proj fusion.2230                 4.8%     512       ~1,750    healthy
input-stall (amortized)               6.9%      —         —        L3: periodic, see strip chart

Artifact 4 — HLO diff: the remat miss and its fix

--- before (bwd computation, excerpt)
+  %rematerialized.ffn.fwd = bf16[8192,28672] fusion(%act_in, %w_up, %w_gate),
+      kind=kLoop, calls=%ffn_fwd_computation      // full FFN fwd re-executed in bwd
+  %custom-call.quant = (bf16[...], f32[...]) custom-call(%rematerialized.ffn.fwd),
+      custom_call_target="fp8_quant_fwd"          // opaque: outputs not in saveable set
--- after (remat policy sees through the op)
-  %get-tuple-element.saved_scales = f32[224] get-tuple-element(%fwd_outputs), index=7
-  // policy: SAVE quant outputs (small: scales f32[224] + int8 acts), recompute only
-  // elementwise epilogue; ffn fwd GEMMs are NOT re-executed
-  attrs: {remat="transparent", pure=true, save_hint="outputs_small"}

Artifact 5 — step-time strip chart (the L3 tell)

Arithmetic check — artifact consistency:
Layer window before: 66 (bwd GEMM) + 25 (recompute) + 45 (exposed RS/AG) ≈ 136 ms; ×80 = 10.9 s + fwd-side tails/opt 2.6 s + amortized stalls 1.04 s + framework 0.5 s ≈ 15.0 s ✓
After: 86 ms × 80 = 6.9 s + tails 0.9 s + framework 0.5 s ≈ 8.3 s ✓ · Recompute total: 25 ms × 80 = 2.0 s ✓ matches fix-(ii) · Exposed: 45 ms × 80 = 3.6 s ✓ matches fix-(i) · Wire under overlap: 36 ms × 80 = 2.9 s < 6.9 s bwd compute → hides ✓

The GPU-fleet twin: what the same investigation prints on H100s

Same incident class on the secondary fleet (256 × H100, 989 TF bf16 → ideal step 3.48 s; observed 13.9 s ≈ 25%). The hierarchy is identical; these are the instruments' actual output formats — quoting them is the difference between "I'd check the kernel summary" and having read one.

Artifact 6 — L0/L2 on GPU: nsys stats --report cuda_gpu_kern_sum report.nsys-rep (top rows)

 Time (%)  Total Time (ns)  Instances   Avg (ns)     Name
 --------  ---------------  ---------  -----------  ----------------------------------------------------------
     21.9    3,050,112,384      2,560    1,191,450  sm90_xmma_gemm_bf16bf16_bf16f32_f32_tn_n_tilesize128x256...
     16.2    2,251,884,096     20,480      109,955  ncclDevKernel_ReduceScatter_Sum_bf16_RING_LL(ncclDevComm...
     13.4    1,866,240,512      2,560      729,000  sm90_xmma_gemm_bf16bf16_bf16f32_f32_tn_n_tilesize128x128...
     13.3    1,852,177,336      2,560      723,500  sm90_xmma_gemm_bf16bf16_... ← SAME shapes as row 1 fwd GEMM,
                                                     appearing in BACKWARD phase = the remat duplicate, GPU edition
      7.7    1,072,693,248     20,480       52,377  ncclDevKernel_AllGather_bf16_RING_LL(...)
      5.8      807,403,520    163,840        4,928  void at::native::vectorized_elementwise_kernel<4, ...>
      4.1      570,425,344      2,560      222,822  flash_bwd_dq_dk_dv_loop_kernel<Flash_bwd_kernel_traits...

Two tells in one table: NCCL kernels at 24% of GPU time (but this table can't say exposed-vs-overlapped — that needs the timeline view: kernel-summary time ≠ critical-path time, the classic nsys misread), and the duplicated GEMM signature in backward — the remat fingerprint without reading any HLO/graph, just kernel names and phase.

Artifact 7 — L2 kernel deep-dive: ncu --set full -k "regex:sm90_xmma.*tilesize128x256" --launch-count 3 (Speed-of-Light excerpt)

  sm90_xmma_gemm_bf16bf16_bf16f32_f32_tn_n_tilesize128x256x64  (8192 x 28672 x 8192)
  Section: GPU Speed Of Light Throughput
    Compute (SM) Throughput           %        54.31      ← vs ~75% healthy for this shape class:
    Memory Throughput                 %        41.02        not starved, not saturated — suspect
    DRAM Throughput                   %        38.77        occupancy/tail, not bandwidth
  Section: Launch Statistics
    Waves Per SM                                 3.42      ← non-integer wave count: 42% tail wave
    Theoretical Occupancy             %        66.67         quantization on 132 SMs at this tile
  Section: Memory Workload Analysis
    L2 Hit Rate                       %        71.88
  ncu comment: est. speedup 12–15% from tile-size re-selection (tail-wave reduction)

The reading discipline: SOL section first (compute% vs memory% names the regime), then wave quantization for the "GEMM is 54% of peak but memory isn't the reason" case — the GPU analogue of TPU tile-padding waste, and the same fix family (shape/tile alignment).

Artifact 8 — L1/L4 on GPU: exposed-NCCL check and the DVFS trap

# L1: timeline-level exposure (nsys GUI or stats nccl + gaps): per-step totals
  nccl wire-time (sum of kernel durations):        1.30 s
  nccl on critical path (MXU-idle overlap gap):    0.85 s   ← exposed: scheduling, and
  backward compute available to hide under:        6.10 s     PyTorch side: check
  → same verdict shape as TPU fix-(i): reorder;      DDP bucket sizes + stream priorities
# L4: before trusting ANY GPU number (the TPU fleet never needs this step):
$ nvidia-smi dmon -s pucm -c 3
  # gpu    pwr  gtemp  mtemp     sm    mem    enc    dec    mclk    pclk
  # 0      688     74     62     99     67      0      0    2619    1830
  # 1      702     81     64     99     71      0      0    2619    1650  ← 180 MHz down:
  # 2      691     76     61     99     69      0      0    2619    1815    thermal throttle on
                                                                            rank 1 = fleet-wide
                                                                            straggler at every
                                                                            synchronous step
$ sudo nvidia-smi -lgc 1830,1830   # lock clocks before ANY comparative measurement

The two GPU-only habits, shown rather than described: kernel-sum NCCL time is not exposure (only the timeline gap is), and clock state is checked before the first measurement is believed — one throttling rank manufactures a "regression" that no software fix will find.

6 · From Incident to Process

Runbook: the §2 audit and L0–L4 hierarchy, verbatim, with per-level tools, exit criteria, and time budgets (audit 0.5 d; L0 attribution 1 d; full attributed decomposition 2 d; no fix begins before attribution reconciles to ±3%). Every incident ends with a one-page report in a shared corpus — searchable by pathology tag — and a runbook delta if any step was missing or wrong. The runbook is versioned and has an owner; a runbook nobody updates is a wiki page, not a process.

Prevention: (1) pre-launch perf review for every new model: geometry checklist (all dims vs tile alignment, with estimated MFU cost per violation — this catches head_dim 96 at config time for free), sharding plan review against the mesh, expected-MFU derived from the roofline calculator + comparable-adjusted (the §2 audit run proactively), and a 30-minute 16-chip pilot profile before the 256-chip run. (2) Regression gates: the model enters the fleet benchmark system (the regression-detection machinery) with its derived expectation as the baseline — "expected 50%" becomes a versioned, derived artifact with error bars, never folklore. (3) MFU-expectation modeling: the roofline calculator is a maintained tool with per-generation machine balances; simulator estimates gate pre-silicon claims.

Organizational design: the ML-perf team owns the runbook and staffs a rotating "perf investigator" seat (also the triage rotation for the regression system — same skills, shared context). Model teams run the pre-launch review with a perf-team reviewer; findings that implicate the compiler are filed with artifacts attached and tracked in a shared queue reviewed with the compiler team biweekly — the queue's age distribution is itself a health metric. Post-incident reviews (blameless, 30 min) update the runbook; twice a year the runbook is exercised as a training drill on a synthetically broken run — the process equivalent of a fire drill, and how L4s become the people who can run this without me.

Staff signal: part (b) answered concretely: entry/exit criteria and time budgets (a process you can't time-box is a vibe), escalation boundaries with the artifact each owner receives, expectations converted from folklore into derived versioned numbers, and a deliberate mechanism for the runbook to stay alive — including training drills. The incident was solved by an engineer; the process makes the next one solvable by whoever is on rotation.

Interviewer follow-ups (§6)
Q: Your pre-launch review adds friction; model teams route around it. How do you keep it real without becoming the perf police?
A: Make the cheap path the reviewed path: the review is 30 minutes with a tool-generated report (geometry check and roofline estimate are automated — the human discussion is only the flagged items), and it comes bundled with things teams want — a derived MFU target they won't be beaten up over, pilot-slice priority, and a named perf contact. Enforcement is economic, not procedural: unreviewed models get no derived expectation, which means any later gap investigation starts with the 0.5-day audit charged to their timeline. Teams route around police; they queue for services.
Q: How do findings actually change the compiler, versus dying in a bug queue?
A: Three mechanisms: every filed bug carries a reproducing HLO + trace + estimated fleet cost in chip-hours (a bug with a dollar number gets triaged differently); the biweekly review is co-owned with a compiler-team TL who commits the queue to their planning, not ours; and patterns that recur ≥3× get promoted from bugs to "heuristic gap" documents that propose the fix shape (like tile-aware fusion caps). The leverage claim to make explicitly: one heuristic fix at the compiler layer is worth N model-team workarounds, so the queue is the highest-ROI artifact the process produces.
Q: What does this process cost, and what's its ROI?
A: Standing cost ≈ one rotation seat plus ~2–4 engineer-weeks per major incident, plus pilot chip-hours (~1% of a training run's budget). Return, from this incident alone: 15.0→8.33 s on a 256-chip run is 44% fewer chip-hours for the same tokens — on a 30-day run that's ~3,400 chip-days recovered, order $1M-scale at internal costs, from roughly three engineer-weeks. The pre-launch review's ROI is better still since it moves the same fixes to before the burn. I'd track exactly this ledger quarterly, because a process that can't state its ROI is the first thing cut.

7 · Remediation Option Space — and When to Stop

Rung (cost ↑)	Levers	When
Config / flags (hours)	sharding specs & `with_sharding_constraint`, overlap scheduling flags, collective bucket sizes, checkpoint cadence/async, input prefetch depth	Always first; also the cheapest confirmation instrument for hypotheses
Compiler-level (days)	layout annotations, fusion/remat hints, XLA flags; file heuristic bugs with repro HLO when the compiler's choice contradicts its own cost model	When the artifact shows a decision the compiler could make better generally — fix the class, not the instance
Model-side (weeks, next revision)	geometry to tile alignment (head_dim 96→128), remat/checkpointing policy, precision plan, sequence/batch shape	When physics says the ceiling moves; costs a quality re-validation, so it rides model-revision trains
Custom kernels — last resort (weeks + forever)	Pallas (TPU) / Triton-CuTe (GPU) for ops the compiler provably can't schedule to the roofline	Only with: profile evidence, roofline headroom proven reachable, and an owner signed up for per-generation maintenance — a hand kernel is a liability re-owned at every compiler and hardware rev

Stopping rule, stated as arithmetic: stop when the next point's cost exceeds its value — here, after 45.0%: remaining levers are a model revision (~2.7 pts, free at next train) and an SDPA Pallas kernel (~1 pt, ~3 engineer-weeks + maintenance). One MFU point on this run ≈ 75 chip-days/month; three engineer-weeks of a perf engineer buys more than that elsewhere in the fleet — the marginal comparison is against the best alternative use of the same engineer, not against zero. Park it, write it down, move the rotation to the next incident.

8 · 60-Minute Delivery Run Sheet

Min	Beat
0–5	Verbatim: "Two answers owed: fix this model, and build the process. First, three framing moves — training or inference, because a single 'inference MFU' is already suspicious — decode is memory-bound by design and MBU is its honest metric; assume training, dense 70B, 256 chips, DP×TP. Second: audit the metric and the expectation before the system — analytic-FLOPs numerator, bf16 dense denominator, stalls included, and I'll show the expectation itself is ~3 points wrong on geometry. Third: the ideal-step arithmetic — 6·70e9·2.1e6 over 256×918 TF is 3.75 s; we're at 15, so 25%, and every hypothesis must cash out in seconds against that line."
5–15	Whiteboard 1: the L0–L4 hierarchy with decision rules (descend by bucket size, exit criteria). Say the L0 split for the incident: 60/24/10/6.
15–38	The worked decomposition, in fix order, doing the arithmetic aloud: exposed-collective (wire 2.9 vs exposed 3.6 → scheduling), remat duplicate (+29% FLOPs, 9 GB free = policy not memory), strip-chart stalls (21 s / 20 steps). Draw the bridge chart; flag the non-linearity of points-in-seconds.
38–48	Process half: runbook (budgets, exits, escalation), pre-launch review (geometry checklist catches head_dim 96 at config time), derived expectations wired into the regression system, the ROI ledger.
48–56	Deep-dives to steer toward if offered: metric-accounting traps · GPU/TPU contrast (DVFS, launch overhead, NCCL, n=10 vs n=3) · when-to-stop economics.
56–60	Close: validation methodology (each fix re-measured, bridge reconciles, residual explained against a corrected target). Two questions back: "What's your current process when a model misses its expectation — and who owns the expectation?" and "How do perf findings flow into the compiler team's planning here?"

Assumptions ledger: chip specs (918 TF, 1.64 TB/s) are published-order figures; effective collective bandwidth (12 GB/s/chip for this pattern), profiler overhead (<2%), and all variance magnitudes are fleet priors to re-measure. Every derived number traces to the §1 arithmetic-check line or a stated per-artifact check.