ReadonlybackwardblockMax vocab-block early-exit epsilon (skips whole cold (token-block, vocab-block) programs, attacking both phases). On at 1e-5 since 2026-07-02; ~lossless when it fires (it skips only blocks whose max prob < eps and which lack the target): alone 1.23× (CPM5) / 2.02× (e4b) at ≤0.004% dh; with the coeff filter at 1e-5/1e-5 the backward is 1.71× (CPM5) / 3.16× (e4b) vs exact. Fires only when the caller supplies blockMax (the production head does); "0" compiles it out.
ReadonlybackwardBackward without steel: the simdgroup_matrix phase-2 kernel (1.41× the scalar kernel, where the H tiling is clean) or the float4 scalar kernel.
ReadonlybackwardApple-CCE coeff filter epsilon (phase-2 per-tile skip). On at 1e-5 since 2026-07-02: the old 0.66%→2.7% dh-accuracy scare was on RANDOM targets (flat softmax, the filter's worst case); the real-data gate (M1 Max, 2026-07-02, chunk-ORPO data) measured CPM5 eps 1e-5 = 0.343% dh for 1.41×, e4b = 0.158% dh for 1.70×, far below the flash-vs-full fp-reassociation scatter that already exists (CPM5 0.850→0.913% relnorm, e4b 1.207→1.220%, cosine ≥0.99993 both). "0" compiles the filter out: exact gradients, the pre-flip kernel.
ReadonlybackwardReadonlybackwardDiagnostics for the scalar backward: skip phase 2's dequantization, or phase 2 altogether (wrong numbers; timing only).
ReadonlybackwardBackward: the steel kernel for both phases wherever its tiling holds.
ReadonlyforwardForward without steel: the simdgroup_matrix kernel (848 ms, the fastest of the three; it needs H % 32 == 0 and otherwise falls to the scalar kernel), the one-thread-per-row scalar kernel (1097 ms), or the lane-sliced kernel (2.6× slower, a recorded dead end; needs whole 32-word rows).
ReadonlyforwardDiagnostic: the simdgroup forward skips its dequantization staging (wrong numbers; timing only).
ReadonlyforwardForward: the steel GEMM kernel (verbatim MLX BlockMMA, 8.4× the old simdgroup forward) wherever its tiling holds; off, the older kernels.
ReadonlyminFewest response rows the flash head serves: a shorter batch takes the exact fused head, which is faster there (the e4b time-tie point, 1024); 0 always takes the flash head.
The flash-CCE head's kernel choices, fixed when the training run is composed (the trainer resolves them once;
FlashCcetakes them at construction) and never read while a kernel dispatches. The defaults are the measured production kernels; every other value is an A/B or diagnostic override.