feat(ck_tile): fp8 sink support for gfx1250 qr_tdm - #11526
Open
steven112163 wants to merge 38 commits into
Open
steven112163 wants to merge 38 commits into
steven112163 wants to merge 38 commits into
Conversation
Adds per-tensor, per-head and per-block granularities. V's descale rides gemm1's E8M0 WMMA scale operand because it varies along that gemm's reduction axis and cannot be recovered afterwards; q and k fold into scale_s and K's per-block descale is applied on the accumulator. P is quantized per 32 elements onto gemm1's A operand.
The fp8 work had reworded or deleted comments that already exist in develop, and added prose that restates the code.
pad_interval, pad_amount, atomic_barrier_address and the iteration config were read unconditionally while still undef.
Packing makes i_nhead a KV head, so one step spans nhead_ratio Q-head tables. Per-head cannot be expressed at all once packed, so exclude it.
seqlen_k was never checked, so ragged lengths silently took the unpadded variant; adding the padded variants also enables group mode. Soft cap and FAST_EXP2=0 have no code path here and now fail to instantiate.
Those scales are constant across gemm1's reduction, so they can ride the output normalization rather than an exponent-only operand that truncates them. BLOCKSCALE keeps the operand and its power-of-two contract. IsSupportedArgument is unrelated to the scales: it revives three checks that were compiled out under NDEBUG.
…x125x only Every suite pinned the scale to TestConfigs<T>::qscale_str, so perhead was never exercised and blockscale only ran from the smoke script at h=1, s=128. Also adds the fp8 and fp8fp32 gtest targets, which had instances but no test at all. Gated on gfx125x so no other target gains test time.
qr_tdm was preferred over qr whenever both matched, including problems whose head dim is mostly padding, where it computes the wrong result.
Take upstream's fp32 softmax reference from #10519: drop our in-place P quantization and let every fp8 reference GEMM1 read the fp32 P, since the device quantizes exp(s-m) and the host would quantize the normalized softmax. Also assert that softmax sink stays off the per-head and block scale granularities: GetSoftmaxScale() has no head index to look their descales up with.
…1003) ## Problem The forward OUT check cannot fail when `ODataType` is `fp8_t`. `check_err` resolves to its fp8 overload, whose fourth parameter is a code distance rather than `rtol`, so the criterion is ``` |out - ref| <= atol OR code_distance <= rtol ``` `get_elimit<FmhaFwdFp8>` returns `atol = 16 * (o_dtype_max > 240 ? 2 : 1)`, i.e. 32 on gfx1250, while O is a convex combination of V and stays within `max|V| = 3`. The first clause is therefore always true. Measured on gfx1250, injecting a fault into the device output immediately before the check: | fault | `-prec=fp8` | `-prec=fp8bf16` | |---|---|---| | none | valid:y | valid:y | | gain x1.05 | **valid:y** | valid:n | | gain x0.72 | **valid:y** | valid:n | | sign inverted | **valid:y** | valid:n | | all zeros | **valid:y** | valid:n | Same binary and configuration for both columns, so this is the fp8-output gate rather than the harness. Tightening the code-distance slot instead does not work: with mixed-sign V about 9% of output elements legitimately change sign under fp8 P quantization, and that overload returns `UINT_MAX` on any sign mismatch. ## Change For the quantized paths, two checks replace the tolerance. **Per element:** `|out - ref| <= u_P * max|V| + u_O * |ref|` `u_P` is taken from `fp8_t` rather than `PDataType`, because the device quantizes P to fp8 whatever the host stores it in; here `PDataType` is float with `quantize_p_ref` applying the per-32 quantization in place, so `numeric_traits<PDataType>` would report 23 mantissa bits and erase the term. Both `u_P` and `u_O` are a full ULP, since the per-32 group scale and the output are each rounded on both sides. `(P|V|)/l` is a P-weighted mean of `|V|`, so `max|V|` bounds it without a second GEMM. **Aggregate:** `alpha = <out - ref, ref> / <ref, ref>` The elementwise bound is structurally blind to a systematic gain: a gain error scales with `|O|` while the bound scales with `max|V| >= |O|`, so where O cancels the threshold rises exactly where the error is. `alpha` is `(g - 1)` exactly for a uniform gain. Non-quantized data types keep the existing `check_err` path. ## Validation, gfx1250, ROCm 10.1 - **180 configurations pass unmodified**, 0 false failures: `fp8`/`fp8bf16`/`fp8fp32` x `n`/`pt`/`ph`/`bs` x hdim 64/128/256 x seqlen 99/256, 128/128, 512/512, 1024/1024, 33/97 x batch and group mode. - **36 of 36 injected faults rejected**: gain 1.05, gain 0.72, sign inversion, zeroed output. - Causal, bottom-right and sliding-window masks pass. - fp16 and bf16 unchanged. Note that ROCm 7.15 miscompiles this path: `no_scale`, `pertensor` and `perhead` return all-NaN there while `blockscale` passes. The numbers above are from ROCm 10.1. ## Note for review `alpha_max = 1e-2` is empirical: the largest `|alpha|` a correct kernel produced across the matrix above was 8.7e-4. Bounding `(P|V|)/l` by `max|V|` is a deliberate simplification; the tight form needs a second host GEMM and buys roughly 2x of elementwise sensitivity, which `alpha` supplies back.
make_v_dram gave qr_tdm a padded V view, so the DMA read past the end of V; with that fixed the head dim narrowing is unnecessary. A new (64,64) tile covers head dim 64 in one k0 loop, and the generator now emits the fp8 elementwise-bias and LSE instances the pipeline already implemented. Two fixes those instances exposed: an all-zero P sub-block rounds to a zero e8m0 scale, which converts to NaN, and the OUT bound applied a normal-number ULP to output that is entirely fp8-subnormal.
Lift the codegen restriction blocking sink generation for fp8/fp8bf16/ fp8fp32 on the gfx1250 qr_tdm pipeline. Fix GetSoftmaxScale() dividing the sink logit by an unfolded scale_s while qr_tdm folds q_descale into its real scale_s for PERHEAD/BLOCKSCALE, which corrupted gpt-oss sink output under those quant modes. Generalize GetSoftmaxScale() so qr_tdm, qr_async_trload, and qr/qr_async share one offset-indexed implementation instead of duplicating the fold logic.
The GPT-OSS sink pre-seeds m before the K loop. Which frame that seed belongs in depends on whether STAGE 2 pre-scales s_acc by scale_s, which happens on the bias paths -- but the predicate tested kHasLogitsSoftCap, which qr_tdm static_asserts to false, so the bias paths silently took the unscaled seed. Every other site in this pipeline classifies the same way the LSE epilogue does: ELEMENTWISE_BIAS and ALIBI carry scale_s in m, the plain path does not. Bring the pre-seed in line, and keep kHasLogitsSoftCap in the disjunction so the classification stays correct if the soft cap is ever implemented here.
Emit qr_tdm instances at head dim 256 for fp16/bf16 and for the whole fp8 family, and give the fp8 dtypes a (256,256) tile to run them on. fp16/bf16 already had the tile; only the pipeline gate was missing. The fp8 family had a commented-out tile whose bk0=64 gave four k0 loops, which check_gemm0_k routes to qr rather than qr_tdm; replace it with bk0 == bk0max == 256 for a single loop, and add the same entry to fp8fp32. bk1 is pinned to 128, the gemm1 warp K step the pipeline static_asserts under kHwGemm1Scale. Both trait sets are narrowed to exclude bias and LSE. qr_tdm wins the runtime dispatch by list order, so an unrestricted gate at 256 would have silently displaced every working qr kernel there. Withholding LSE keeps the def_lse=true d=256 suites on qr_ks_vs, so qr stays exercised at this head dim. fp8 has no such safety net -- its qr list only takes tiles with k0_loops >= 2, so at 256 qr_tdm is the sole provider.
The prefill mainloop applied the sink->normal window jump to K, V and bias together at i_total_loops == num_sink_loop - 1, but the three are not consumed on the same schedule. K ping-pongs through ptrk0/ptrk1 and is prefetched two iterations ahead of consumption; V and bias are only one ahead. Jumping all three at the same iteration left K reading the first normal-region tile from the pre-jump address. Split the jump by consumer: - K fires at num_sink_loop - 2. When num_sink_loop == 1 no such iteration exists, so the prologue's second prefetch applies it instead -- that prefetch already targets what the mainloop consumes at iteration 0. - V moves before its own prefetch at the top of the loop body rather than after, since that prefetch feeds the following iteration. - bias keeps the original gate, which matches its 1-ahead schedule. Cherry-picked from the qr_tdm portion of #10616, which fixes the same class of bug across qr, qr_async and qr_tdm.
Bound the KV descale index by the real seqlen_k, since the mask's tile range rounds past it, and divide the softmax sink by the same folded scale_s the pipeline re-applies so per-head and per-block descales cancel.
Emit -Rpass-analysis=kernel-resource-usage remarks for the FMHA fwd and bwd instance targets, giving per-kernel VGPR/SGPR/LDS/scratch/spill counts and the occupancy ceiling without running the kernel. The flag is attached to the instance targets rather than passed through CMAKE_CXX_FLAGS because the generated blobs are compiled as LANGUAGE HIP and therefore only see CMAKE_HIP_FLAGS.
IsOutOfSinkBound short-circuits four divergent per-element comparisons with &&/||. Each one makes the compiler emit an exec-mask save, a branch and a restore, which the sink kernels pay on every masked tile. Replace the short-circuits with bitwise & and | under __gfx1250__. On the gfx1250 ASM dump this removes all 192 extra exec-mask sites across the 32 masked sink kernels, cuts the mean instruction delta from +1492 to +264 and the branch delta from +65 to +1. Measured sink cost against the matching baseline drops from +3.66% mean / +10.01% worst to +1.34% / +3.98%. bool locals are needed because the return type is deduced and bool & bool yields int.
github-actions
Bot
added
project: composablekernel
ck: attention-moe
A descale is dequantized_max/quantized_max, but quantized_max was taken from the data type rather than from what the fill actually wrote. The two agree only for init=3, which fills to the fp8 maximum. Every other fill leaves headroom, and anchoring on the format shrank the logits by the ratio between them until the softmax was uniform to within a rounding step, at which point one shared fp8 rounding decision on P became a uniform gain on the whole row. - Measure the fill magnitude instead of using numeric<T>::max() - Apply to pertensor, blockscale and perhead, which share the anchor - Drop the init=3 special case on the blockscale V exponent, now subsumed
The sink and the quantization scales were only ever exercised apart, so the case that matters was untested: a sink pre-seeds the running max above the true row max, which makes max P arbitrary rather than exactly 1, and only the derived check the descale paths use can see the uniform gain that follows. The stock tolerance cannot. Sweep the axis on QuantScale, which already covers all four scales, rather than on SinkWindowMask, which runs one fixed scale per config and so would never cross the two. - gpt-oss: one learned per-head logit, keeps the shape's own mask - StreamingLLM: an always-attended KV prefix, which needs a window mask and so replaces the shape's mask
get_elimit<FmhaFwdFp8> passed (0, 16|32) into check_err's fp8 overload, whose third parameter is a count of e4m3 grid steps rather than rtol. That made the step clause unsatisfiable and the absolute clause always true, so the check could not fail: the gtest's qscale=n cells still passed with OUT scaled by 2. Pass the step count in the slot that means step count, and express atol relative to max|ref| since a correct OUT deviates in proportion to the tensor's own range and that range spans three orders of magnitude across the init methods.
alpha averages the query rows that hold an unmasked key, because the fp8 rounding of P is shared across a row, so its noise floor is per-shape and a single 1e-2 fits neither end of the matrix.
Only the qr_tdm V scale rides an E8M0 operand, so keep upstream's random descale fill and give v its own E8M0-exact fill on gfx1250. Snapping after a uniform fill collapsed every scale onto a single exponent.
…nto users/steven112163/fmha-gfx1250-tdm-sink # Conflicts: # projects/composablekernel/example/ck_tile/01_fmha/fmha_fwd_runner.hpp # projects/composablekernel/include/ck_tile/ops/fmha/kernel/fmha_fwd_kernel.hpp
Per-head descales now offset by head and by batch in both modes, the reference varies along both axes so a wrong offset shows up, and unsupported arguments return instead of throwing.
A descale is dequantized_max/dtype_max, which treats the input tensors as the image of an fp32 tensor whose amax is qkv_max. That only holds when the fill saturates the format, and init=3 alone does. Under any other fill the descale contradicts its own tensor: at init=uf the logits land near 1e-4, a sink token becomes the row maximum, every P value collapses onto a single number, and its rounding turns into a pure gain on the output. Reject the combination rather than re-deriving the descale from the data, which produced plausible output and hid the misuse. Also restore the blockscale descale fill to a narrow uniform band. The pipeline folds k_descale into s_acc per N-block, so a wide power-of-two spread there peaks every row onto whichever KV block carries the largest scale. That drove the effective key count to the corner the OUT systematic gain bound is calibrated on and left it no margin.
…fmha-gfx1250-tdm-sink
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
kv_tile_start indexed the per-block K and V descale arrays as kv_load_start + i_total_loops * kN0, but k_origin, which picks the tile the kernel actually loads, is sink-aware. When the StreamLLM sink phase is non-empty kv_load_start is forced to 0, so during the normal phase the descale index trailed the loaded column by physical_seqlen_k_start - sink_seq_end and every KV block past the sink->normal jump was scaled by a descale belonging to an earlier block. The wrong K descale multiplies s_acc before the softmax, so the error lands inside the exponent: measured worst 33.8x on 87120 elements at d128, up to 62.2x elsewhere. Derive kv_tile_start with the same branch k_origin used and let k_origin and the ALIBI current_k_origin reuse it, in both mainloops. That removes four hand-copied lambdas, so the descale index and the loaded column can no longer drift apart. Only blockscale is affected. Pertensor and perhead fold both descales into one scale_s scalar with no per-block index, and no_scale has none. This is the case static_assert(!(kBlockScale && kHasSink)) guarded before it was removed. The removal was too broad rather than wrong: the sink phase is non-empty only when GetSinkTileRangeAlongX takes its else branch, which needs x_start > sink_seq_end, and x_start is nonzero only under IsLocal. kStreamLlmMask is causal, so the QuantScale sweep always collapsed the phase to zero tiles and never reached this path. Add a local-window tuple to QuantScale to cover it. Verified on gfx1250: the descale sweep goes 21 failures to 0 across the same 33 shapes, and the quantization-scale suite is 288/288 on fp8, fp8bf16 and fp8fp32.
Revert both halves of "route hdim 256 to qr_tdm on gfx1250": the d256 pipeline gate and the fp8 (256,256) tile in the codegen, and the d256 head-dim axis in the forward tests. Head dims above 128 belong with the wider large-hdim work rather than with the sink fixes this branch carries. That axis was also the only source of failures in the forward suite. fp16 QuantScale reported 16 wrong cases at d256 and none at d128, every one of them 1 to 4 elements at one or two fp16 ULP. bf16 compiles an identical instantiation and passed, its relative tolerance being about eight times looser, which places the problem in the d256 tolerance rather than in the sink path. Four of the 16 came from tuples that predate this branch, so the bound needs revisiting with the large-hdim work regardless. Verified on gfx1250 after the revert: the fp8 quantization-scale suite is 168/168 on fp8, fp8bf16 and fp8fp32, and the fp8 qscale smoke matrix is 48/48 at d128 with no missing instance.
The local-window tuple added with the sink-aware descale index uses sink=128, which equals kN0 and so gives num_sink_loop 1. That reaches the prologue jump alone; the mainloop branch at i_total_loops == num_sink_loop - 2 stays unvisited, and it is the branch the descale index depends on most. Add sink=512 for num_sink_loop 4, with GQA ratio 2 so the descale stride is crossed as well. Also record on kStreamLlmMask why the causal row cannot carry a live sink phase: left=-1 makes y == y_total, so x_start is 0 and GetSinkTileRangeAlongX collapses the phase to zero tiles, leaving that row as a check that the has_sink instantiation still matches the sinkless answer. Verified on gfx1250: 168/168 on fp8, fp8bf16 and fp8fp32.
steven112163
changed the title
steven112163
marked this pull request as ready for review
September 15, 2026 05:02
The Static checks CI stage runs clang-format-18 across the tree. Three files this branch touches were never formatted, so the stage failed and every downstream build and test stage was skipped. Reflow only: rejoined || operands, rewrapped the QuantScale base clause, reflowed the fmha_fwd_run argument list, and fixed comment indentation inside the CK_TILE_FMHA_FORCE_HEAD_MAJOR guard. No tokens changed.
LJ-underdog
reviewed
Sep 18, 2026
The sink pre-seed in qr_ks_vs_tdm has to match the unit convention of the
branch that consumes m: log2-domain under bias or alibi, raw-qk without it.
Only the elementwise-bias row exercised that, leaving alibi with no coverage
anywhere.
On gfx1250 the dispatcher sends fp16/bf16 hdim128 + {bias, alibi} + sink to
qr_tdm, so the new row lands there. fp8 has no alibi instance and skips it.
Also drop the claim that no fp8 pipeline is generated with bias. The gfx1250
codegen emits bias for fp8, and fp8 bias + sink is carried by SinkWindowMask.
Every descale is dequantized_max/quantized_max with quantized_max taken as the format maximum, so it only describes a fill that saturates the format, and init=3 alone does. A validated run under any other fill reads numbers that contradict their own tensor. A timing run reads none of them: do_validation == 0 skips verification entirely, yet the guard still refused to start, so -qscale=pt -init=uf could not be profiled. Gate on do_validation, matching the p_drop precedent, and move the check up beside the other argument checks so it fires before the host buffers and the alibi slopes are allocated.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
5 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Motivation
qr_tdmis the only pipeline gfx1250 dispatches for fp8 attention, but sinksupport was blocked there by codegen and by three defects that only surface once
a sink region is present. gpt-oss sink output was corrupted under PERHEAD and
BLOCKSCALE quantization, the bias paths took an unscaled pre-seed, and the
per-block descale index drifted away from the tile actually loaded whenever the
StreamLLM sink phase was non-empty. The sink-bound predicate also cost four
divergent branches per tile in the mainloop.
JIRA ID: ROCM-30734, AICK-2259
Technical Details
fp8/fp8bf16/fp8fp32 on
qr_tdm, and fixGetSoftmaxScale()dividing the sinklogit by an unfolded
scale_swhile the pipeline foldsq_descaleinto itsreal
scale_sfor PERHEAD/BLOCKSCALE, which corrupted gpt-oss sink outputunder those quant modes.
kHasLogitsSoftCap, whichqr_tdmstatic_asserts to false, so the bias paths silently took the unscaledseed. Classify on bias, matching every other site in the pipeline.
V and bias only one. Split the sink->normal jump by consumer instead of firing
all three at the same iteration.
kv_tile_startindexed the per-block K and Vdescale arrays as
kv_load_start + i_total_loops * kN0, butk_origin, whichpicks the tile actually loaded, is sink-aware. With a non-empty StreamLLM sink
phase
kv_load_startis forced to 0, so every KV block past the sink->normaljump was scaled by a descale belonging to an earlier block. The wrong K descale
multiplies
s_accbefore the softmax, so the error lands inside the exponent.Both mainloops now derive the index from the same branch
k_originuses. Onlyblockscale is affected; pertensor and perhead fold both descales into one
scale_sscalar with no per-block index, and no_scale has none.IsOutOfSinkBoundshort-circuited four divergentcomparisons, each emitting an exec-mask save/branch/restore. Bitwise form under
__gfx1250__.the fp8 quantization-scale matrix, and add local-window tuples that make the
sink phase non-empty.
sink=128 == kN0givesnum_sink_loop1 and reaches theprologue jump;
sink=512givesnum_sink_loop4 and reaches the mainloop jumpat
i_total_loops == num_sink_loop - 2, which the descale index depends onmost. The causal
kStreamLlmMaskrow cannot carry a live sink phase at all,since
left=-1makesy == y_totaland collapses the phase to zero tiles.hdim 256 routing was part of this branch and has been reverted; it moves to the
large-hdim work (256, 512) together. It was also the only source of failures in
the forward suite: fp16
QuantScalereported 16 wrong cases at d256 and none atd128, each 1 to 4 elements at one or two fp16 ULP. bf16 compiles an identical
instantiation and passed, its relative tolerance being roughly eight times
looser, which places the problem in the d256 tolerance rather than in the sink
path. Four of those 16 came from tuples that predate this branch.
Test Plan
The descale-index defect is isolated with a 33-run matrix that varies mask kind,
sink size, seqlen, head dim and dtype around the failing shape, each run
predicted pass or fail in advance from the index arithmetic. A control build and
a fixed build are run over the same matrix, so the fix is judged by whether the
predicted set flips, not by a count alone.
Coverage for the fix lives in
TestCkTileFmhaFwd/QuantScale, which crosses thefour quantization scales with the three sink kinds. Two local-window tuples carry
the live sink phase:
sink=128for the prologue jump andsink=512for themainloop jump at
num_sink_loop - 2. Causal rows cannot reach either, since thesink phase collapses to zero tiles.
The fp8 qscale smoke matrix checks that every
(precision, scale, mode, perm)combination still has an instance and validates, separating a missing instance
from a numerical failure rather than collapsing both into one exit code.
Finally the ROCM-30734 repro re-runs the 15 GTest cases the ticket reports, in
fp16 and bf16, matching each index against the shape it is supposed to be so a
renumbered
Combine()product is reported rather than silently counted.The sink pre-seed gate is isolated the same way: the branch is built twice with
nothing different but the two pre-seed conditions, the decode path and the
prefill copy. The control reverts them to
if constexpr(kHasLogitsSoftCap),which
qr_tdmstatic_asserts to false, so every path takes the raw-qk seed. Thefixed build is this branch as it stands. Both run the whole
SinkWindowMasksuite on fp16, bf16 and fp8, so 24 cases per dtype are compared across a
one-condition diff.
Only the
init_sink=1rows can discriminate. That fill draws the sink uniformfrom [-1,1], so a positive value exists and an inflated seed takes over the row
max;
init_sink=0leaves the sink at zero, where both unit conventions agree andthe defect is invisible by construction. The rows carrying the signal are the
elementwise-bias and alibi tuples at
{128, ..., 1024, "t:128,30", 1}.That
qr_tdmis the pipeline under test is checked rather than assumed:generating the dispatch chain for gfx1250 shows fp16/bf16 hdim128 +
{bias, alibi} + sink and fp8 hdim128 + bias + sink all resolving to
QRKSVS_TDM, which occupies all four pad variants ahead ofqr. fp8 has noalibi instance, so that row skips there.
Test Result
Descale index, control build -- 21 of 33 fail, exactly the predicted set
Worst observed error at d128: 33.84x on 87120 elements (fp8), 44.49x on 108813
(fp8bf16), 45.83x on 110625 (fp8fp32).
Descale index, fixed build -- 33 of 33 pass
Sink pre-seed gate -- 10 failures to 0 across fp16, bf16, fp8
Control build,
SinkWindowMaskon gfx1250:Fixed build, same suite and same shapes:
Every control failure is a bias or alibi row with
init_sink=1, in both batchand group mode. The
init_sink=0bias row passes in both builds, as does everyrow without bias. The alibi tuple is new in 03be482; alibi with a sink had no
coverage anywhere before, and it fails on fp16 and bf16 without the fix.
fp8 quantization-scale unit suite -- 168/168 on each dtype
fp8 qscale smoke matrix, d128 -- 48/48, no missing instance
ROCM-30734 repro, fp16 + bf16 -- 30/30, both clusters fixed
The 12
ElementwiseBiascases went from 52 to 97 wrong elements down to 0, andthe
SinkWindowMaskcase that was 131072 wrong (50%) is 0. Every case stillreports a
qr_tdm+...+sink+instance, so the fixes are exercised on the intendedpipeline rather than dispatched elsewhere. Index-to-shape checks all matched, so
no result is a renumbering artifact.
Branch-free sink bound -- exec-mask sites and sink cost
The branch-free rewrite removes all 192 extra exec-mask sites across the 32
masked sink kernels; mean instruction delta drops from +1492 to +264 and branch
delta from +65 to +1. Sink cost against the matching baseline goes from
+3.66% mean / +10.01% worst to +1.34% / +3.98%.
Measured on an earlier revision of this branch, before the descale-index fix and
the hdim 256 revert. The kernel set has changed under those figures, so the
counts above include d256 variants this PR no longer ships; the per-kernel deltas
do not depend on head dim.