CUDA: MMQ tile configuration and FP4 accumulation for GB10 (sm_121) - #198
CUDA: MMQ tile configuration and FP4 accumulation for GB10 (sm_121)#198danielhanchen wants to merge 5 commits into
Conversation
Dense prompt matmuls on GB10 run the MMQ mainloop at one block per SM: the 128-row tile needs 61.5 KB of shared memory at J=128 and the kernel asks for 256 threads, so nothing can overlap the register-staged global loads or the two barriers per K step. Halving the tile to 64 rows with 128 threads and occupancy 2 makes two blocks resident and is worth 9 to 13 percent on the K-quants at prompt-sized batches. mul_mat_id does not want this: each expert only holds ncols_max*n_expert_used/n_experts columns, so halving the tile height only doubles the number of row tiles that pay for a mostly empty J tile. The tile table therefore keeps the stock 128-row tile in its fallback == true entries and mul_mat_id asks for that config on Blackwell; requesting the bounds-checked variant cannot change results, it only enables the src0->ne[1] range check. The block-scaled FP4 MMA already takes C as an accumulator input, but the vec_dot zeroed a temporary tile per instruction and added it to the running sum afterwards. Accumulating into the sum directly removes 805 FADD and all of the register spilling from the NVFP4 J=128 kernel. Both changes are confined to the Blackwell configuration and to code under BLACKWELL_MMA_AVAILABLE, so no other GPU sees a different kernel.
…re-deriving it mul_mat_q_case re-derived the fallback flag from args.nrows_x % 128, so the flag that ggml_cuda_mul_mat_q computes (and that already sizes the src1 padding through ggml_cuda_mmq_get_J_max) never reached the kernel selection. It matched by construction while the only input was ne01, but any other reason to want the bounds-checked configuration was silently dropped, including the mul_mat_id case added in the previous commit. Carry the decision in mmq_args so the configuration the padding was sized for is the configuration that runs. No behaviour change for the ne01-only rule.
|
@codex security review |
Codex Review SummaryThis comment shows the latest Codex review activity on this pull request.
ℹ️ About Codex in GitHubYour team has set up Codex to review pull requests in this repo. Reviews are triggered when you
Codex reacts with 👀 while any review is running, comments if it has suggestions, and reacts with 👍 once all reviews finish with no findings. |
|
Codex Review: Didn't find any major issues. Another round soon, please! Reviewed commit: ℹ️ About Codex in GitHubYour team has set up Codex to review pull requests in this repo. Reviews are triggered when you
If Codex has suggestions, it will comment; otherwise it will react with 👍. Codex can also answer questions or update the PR. Try commenting "@codex address that feedback". |
What this changes, in four questions1. Before the PR, what happened? On GB10 (sm_121) the Blackwell MMQ tile table used the stock tile configuration for the K-quants and the FP4 types, and the block-scaled MMA path re-derived its accumulator handling rather than accumulating into the running sum. 2. After the PR, what happens? The tile table carries a 64-row / occupancy-2 tile for those types, which is a win for dense 3. Is this a real issue or a fake one? Real, and it is a throughput issue rather than a correctness one, which is worth stating plainly. Any user running a K-quant or FP4 GGUF on a GB10 gets the tile shape chosen for a different Blackwell part. The 4. If merged, does it break anything? This is the one PR in my batch that is not RPC-gated, so it is live for every CUDA user, and that is the bar it was held to:
Verified: builds clean with Also carries a comment-reduction pass, AST-gated as comments-only against the pre-pass head. |
|
One coverage note that bears directly on the gate here. The dispatch predicate is static bool blackwell_mma_available(const int cc) {
return GGML_CUDA_CC_IS_NVIDIA(cc) && ggml_cuda_highest_compiled_arch(cc) >= GGML_CUDA_CC_BLACKWELL &&
ggml_cuda_highest_compiled_arch(cc) < GGML_CUDA_CC_RUBIN;
}Note it tests What changes for sm_120: Q4_K, Q5_K and Q6_K previously had no Blackwell entry and fell through to the Ampere config const bool fallback = ne01 % 128 != 0 || (ids && blackwell_mma_available(cc));Q4_K covers essentially every Q4_K_M GGUF, so this is the common case on a 5090. This is measured, not inferred. Compiling every MMQ translation unit (
Correctness on GB10 is fine: I could not measure the sm_120 performance impact, because there is no RTX 50-series GPU on the pair I tested on. So the position is narrow and specific: the kernels provably change for that population, and the effect on them is unmeasured.
Narrowing the dispatch to an exact-equality check on |
The swept tiles were written into mmq-config-blackwell.cuh, so every sm_12x device got a table measured on one DGX Spark. GB10 is one SM cluster against unified LPDDR5X; a discrete Blackwell part has neither, and nothing in the sweep was measured on one. The tiles move to mmq-config-dgx-spark.cuh and mmq-config-blackwell.cuh returns to master byte for byte. Types the sweep did not cover fall through to the Blackwell table, so this narrows the blast radius without narrowing coverage. Dispatch is an exact equality against GGML_CUDA_CC_DGX_SPARK, not a >= range, following mmvq.cu:111, which tests the same ggml_cuda_highest_compiled_arch the existing dispatch uses. fattn.cu:194 is a >= threshold and is deliberately not the pattern copied here: a range would catch every future sm_12x part with a table measured on none of them. The host and device selectors mirror each other, or the host would size shared memory for one config while the kernel compiled another. Verified by SASS rather than by reading. Compiling the q4_k, q5_k, q6_k and q8_0 MMQ instantiations for sm_75, sm_80, sm_90, sm_120 and sm_121a and disassembling each cubin, against UPSTREAM MASTER as the reference: sm_75, sm_80, sm_90, sm_120: match master for all four types sm_121a: differs for q4_k, q5_k, q6_k; matches master for q8_0 which is exactly the intended shape, since q8_0 is not in the Spark table. The residual on every matching cell is 32 lines carrying one constant, 0x3bf to 0x3ca, an assert's __LINE__ shifted by the 11 lines this adds above it, materialised as MOV on sm_75, IMAD.MOV.U32 on sm_90 and HFMA2 on sm_80 and sm_120. No instruction, register allocation or scheduling differs, so the claim is "unchanged apart from an assert's line number" rather than byte-identical. ggml-cuda builds green for sm_121a.
|
@codex review |
|
Codex Review: Didn't find any major issues. What shall we delve into next? Reviewed commit: ℹ️ About Codex in GitHubYour team has set up Codex to review pull requests in this repo. Reviews are triggered when you
If Codex has suggestions, it will comment; otherwise it will react with 👍. Codex can also answer questions or update the PR. Try commenting "@codex address that feedback". |
Simulation against the merge base,
|
| arch | 585f9c9eb (before the narrowing) |
d0a262e1a (after) |
|---|---|---|
| sm_75 | 0 of 23 | 0 of 23 |
| sm_80 | 0 of 23 | 0 of 23 |
| sm_90 | 0 of 23 | 0 of 23 |
| sm_120a | 5 of 23 | 2 of 23 |
| sm_121a | 5 of 23 | 5 of 23 |
Before the narrowing, sm_120a and sm_121a got identical SASS in all five units (147760, 136048,
118597, 123954, 135110 in both columns), i.e. an RTX 50-series card was running a GB10-swept tile
table verbatim. That was worth fixing.
Corroborated by a 13200-row host selector matrix (31 compute capabilities x 20 types x 11 J x both
fallback values, evaluating the real ggml_cuda_mmq_get_config):
- with sm_121 in the arch list:
585f9c9ebchanges 220 rows on cc 1200/1201/1205/1209,d0a262e1a
changes 0; - with an arch list that does not contain sm_121 (
750,800,900,1200, a consumer-Blackwell CI
build):585f9c9ebchanges 550 rows,d0a262e1achanges 0 of 13200; - all six AMD encodings: 0 rows in either tree.
Finding 1: the narrowing does not cover the FP4 accumulator change
d0a262e1a narrows the tile selector. The other half of the PR, the change in
ggml/src/ggml-cuda/mmq-vec-dot.cuh, is inside ggml_cuda_mmq_vec_dot_fp4_fp4_mma, compiled for
every BLACKWELL_MMA_AVAILABLE arch, which is __CUDA_ARCH__ in [1200, 1300), all of sm_120 too.
That is why sm_120a is 2 and not 0: mmq-instance-mxfp4 (139680 SASS lines) and
mmq-instance-nvfp4 (127968).
It is bit-neutral, and that was measured, not assumed. A tree consisting of the merge base plus
that one hunk and nothing else, against the merge base, GGML_CUDA_FORCE_MMQ=1, 8 types x 6 shapes
including m=255 to force the fallback config, FNV-1a over the raw f32 output, every case computed
twice to prove determinism and the whole set reproduced from a second invocation:
| cases differing | |
|---|---|
| base vs base + FP4 accumulator hunk, all 48 | 0 |
PTX says of mma that "The accumulation order, rounding and handling of subnormal inputs are
unspecified", so chaining C through the block-scaled MMA instead of zeroing a temporary could have
changed results. It does not.
So sm_120 gets an unmeasured code-generation change, not an unmeasured numerics change. Still, the
header comment in mmq-config-dgx-spark.cuh says "nothing here was measured on a discrete
Blackwell part" and "the override is exactly the swept entries", and that is not true of the diff
as a whole while this hunk is ungated. Wrapping it in
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ == GGML_CUDA_CC_DGX_SPARK takes sm_120a from 2 to 0.
Finding 2: the tile change is a numerics change on GB10
base vs head, same harness, sm_121a:
| type | n=1 | n=16 | n=128 | k=1024 m=512 n=512 | k=2048 m=1024 n=2048 | m=255 (fallback) |
|---|---|---|---|---|---|---|
| nvfp4 | same | same | same | DIFFERS | DIFFERS | same |
| mxfp4 | same | same | same | same | same | same |
| q4_K | same | same | same | DIFFERS | DIFFERS | same |
| q5_K | same | same | same | DIFFERS | DIFFERS | same |
| q6_K | same | same | same | DIFFERS | DIFFERS | same |
| q4_0, q8_0, iq4_xs | same | same | same | same | same | same |
8 of 48. Only the five types the table names, only where the 64-row tile changes the stream-k
split, and the m=255 fallback column is identical for every type, which is the table falling
through to the Blackwell entries as intended.
Greedy generation follows from it. llama-cli -ngl 99 --temp 0 --top-k 1 --seed 42 -n 160 --single-turn, Qwen3.5-4B-UD-Q4_K_XL, single device, default clocks, no --rpc:
| run | md5 of generated text |
|---|---|
| base, run A | ebaa06f1579d5edd |
| base, run B | ebaa06f1579d5edd |
| head | e404e817a0e8f7c2 |
The base/base control is byte-identical, so the divergence is real and not harness noise. Output
splits mid-sentence. This is a floating point reassociation, not a correctness bug, and it is
confined to GB10: sm_75/80/90 are byte-identical in SASS and sm_120's only change is bit-neutral.
But the PR body's "Same kernels, same numerics" is not right for a change that alters I and
nthreads under stream_k, because the tile shape decides how K is split across CTAs and
therefore the order the partial sums are combined. Worth correcting in the body, since a reader
will otherwise assume a Spark deployment can take this without revalidating output.
Everything else is unchanged
test-backend-ops -b CUDA0, single device, no--rpc:-o MUL_MAT1193/1193 on base and
1193/1193 on head;-o MUL_MAT_ID869/869 and 869/869. Counts match base exactly.
MUL_MAT_IDpassing matters: it is the operation the new|| (ids && spark)term changes.- Shared memory: GB10 reports
MAX_SHARED_MEMORY_PER_BLOCK_OPTIN101376 B and 102400 B per SM. Over
the 55 rows the table changes, the base config asks 57856 B at occupancy 1 (a second block would
need 115712 B) and the new one 38400 B at occupancy 2 for 76800 B. Nothing in the 13200-row matrix
exceeds 101376 B on any tree. fallback:mmq_argshas exactly two construction sites, both inggml_cuda_mul_mat_q, both
updated, so nothing leaves the new field unset. At cc=1210 everyfallback == truerow is
byte-identical to the merge base, and the Blackwell table'sfallbackrows differ from its
non-fallback rows only in the flag, so forcing it formul_mat_idadds the range check and
nothing else, exactly as the body says.- HIP: the device branch is inside the non-HIP arm of
#ifdef GGML_USE_HIPand additionally tests
defined(__CUDA_ARCH__); the host branch sits after the unconditionalGGML_CUDA_CC_IS_AMD
return and is itself gated onGGML_CUDA_CC_IS_NVIDIA(cc). There is no AMD hardware on this
pair, so this is reasoned from source and not tested. The host half is executed evidence: 0 of
1320 AMD rows differ. - One caveat on "exact equality": the test is exact on
ggml_cuda_highest_compiled_arch(cc), not on
cc, so any cc from 1210 up to the next compiled arch also selects the table. No such part
exists, and the merge base has a skew of the same shape for the same hypothetical cc, so this is
neither introduced nor fixed here.
Thermal guard log had 559 lines before and 559 after the measured window, so no cap overlaps any
number above.
The gate narrowing in d0a262e was entirely in the tile configuration selector. The other half of this PR, the FP4 accumulator change, lives in ggml_cuda_mmq_vec_dot_fp4_fp4_mma, which is compiled for every architecture defining BLACKWELL_MMA_AVAILABLE, i.e. __CUDA_ARCH__ in [1200, 1300). That is all of sm_120 as well as sm_121, so an RTX 50-series card was still getting it. Measured with cuobjdump -sass over all 23 MMQ translation units at 75-real;80-real;90-real;120a-real;121a-real, against the merge base, with each differing line classified rather than filtered (the +11 __LINE__ shift this PR introduces in mmq.cuh is compared on the instruction encoding word, since ptxas renders the same constant as MOV on sm_75 and HFMA2.MMA on sm_80): arch before the narrowing at d0a262e with this commit sm_75 0 of 23 0 of 23 0 of 23 sm_80 0 of 23 0 of 23 0 of 23 sm_90 0 of 23 0 of 23 0 of 23 sm_120a 5 of 23 2 of 23 0 of 23 sm_121a 5 of 23 5 of 23 5 of 23 The two that remained on sm_120a were mmq-instance-mxfp4 (139680 SASS lines) and mmq-instance-nvfp4 (127968). GB10 loses nothing. A strict comparison of d0a262e against this commit is byte-identical on sm_121a for all 23 translation units, so the Spark runs the same machine code either way. The change being gated is bit-neutral, so this is about not shipping unmeasured code generation rather than about results: the merge base plus that hunk alone, against the merge base, is bit-identical on all 48 (type, shape) MUL_MAT cases including MXFP4 and NVFP4. It is gated anyway, because PTX leaves mma's accumulation order unspecified and the hunk was swept on a DGX Spark and nowhere else, which is the same reason the tile table is gated.
|
Pushed The commit gates the FP4 accumulator hunk in
GB10 loses nothing: a strict SASS comparison of The body now says what was measured instead of "same kernels, same numerics": the 8-of-48 bitwise
|
1 participant
Prompt processing on GB10 (DGX Spark, sm_121a) is matmul bound, and MMQ is the kernel that runs every
quantized
mul_mat. An audit of the sm_121a SASS put the MMQ mainloop at about 31 % of the int8tensor-core peak and 20 to 27 % of the FP4 peak at prompt-sized batches, against 65 to 71 % for
CUTLASS on the same GPU. Two of the causes are cheap to remove, and both are confined to the DGX
Spark configuration: an exact
== GGML_CUDA_CC_DGX_SPARKon the host and the same test on__CUDA_ARCH__in device code, for the tile table and for the FP4 accumulator alike.1. The dense tile is one block per SM. At J = 128 the 128-row tile needs 61.5 KB of shared memory
and the kernel asks for 256 threads, so exactly one block is resident and nothing overlaps the
register-staged global loads or the two
__syncthreadsper K step. The MMA layout ties I to 16 rowsper warp, so a 64-row tile at 128 threads is the only way to get two resident blocks (40.5 KB per
block at J = 128, two fit in the 100 KB per SM of GB10). Same kernels and more resident warps.
Not the same numerics on GB10, and that correction matters to anyone deploying this. An earlier
version of this section said "same numerics". It is wrong. These kernels use stream-k, which splits
the K dimension across CTAs and combines the partial sums afterwards, and the tile shape is what
decides that split. Changing
Ifrom 128 to 64 therefore reassociates a floating point sum.Measured directly, one
MUL_MATper (type, shape) on the CUDA backend withGGML_CUDA_FORCE_MMQ=1and a hash of the raw f32 output, 8 types x 6 shapes, every case computed twice to confirm
determinism:
8 of 48, all on the five types this table re-tiles, all where the 64-row tile changes the stream-k
split. The types deliberately left on the stock tile are bit-identical everywhere, and the
m=255fallback column is bit-identical for every type.
Greedy generation follows from it.
llama-cli -ngl 99 --temp 0 --top-k 1 --seed 42 -n 160 --single-turnon a Q4_K model, single device, default clocks, no--rpc:ebaa06f1579d5eddebaa06f1579d5edde404e817a0e8f7c2The base-against-base control is byte-identical, so the divergence is the change and not
harness noise. The output splits mid-sentence.
This is not a correctness bug.
test-backend-ops test -b CUDA0gives-o MUL_MAT1193/1193 and-o MUL_MAT_ID869/869, matching the merge base exactly on both. A reassociated FP32 sum is whatevery MMQ tile change in this file does, and PTX states that for
mma"The accumulation order,rounding and handling of subnormal inputs are unspecified".
It does mean a DGX Spark deployment cannot take this without revalidating output. Anything
pinned to an exact generation, a cached response, a golden transcript or a reproducibility claim on
GB10 will move. No other GPU is affected: sm_75, sm_80 and sm_90 are byte-identical in SASS to the
merge base, and sm_120's only remaining difference is bit-neutral (see below).
mul_mat_iddoes not want this. The MoE launch picks J fromncols_max= the total token countwhile each expert only holds
ncols_max*n_expert_used/n_expertscolumns, so halving the tile heightonly doubles the number of row tiles that pay for a mostly empty J tile: measured at -12 to -21 % on
MUL_MAT_IDat 2048 tokens. The table therefore keeps the stock 128-row tile in itsfallback == trueentries andmul_mat_idasks for that configuration on Blackwell. Requesting thebounds-checked variant cannot change results, it only enables the
src0->ne[1]range check in thetile loads and the write-back.
2. The FP4 accumulator is copied.
mma.sync...m16n8k64block-scaled already takes C as anaccumulator input, but
ggml_cuda_mmq_vec_dot_fp4_fp4_mmazeroed a temporary tile per instructionand added it into
sum[]afterwards. Accumulating intosum[]directly removes 805 FADD and all ofthe register spilling (26 STL / 30 LDL) from the NVFP4 J = 128 kernel; 224 -> 256 OMMA visible in the
loop.
The second commit fixes a latent bug found on the way:
mul_mat_q_casere-derivedfallbackfromargs.nrows_x % 128and ignored the flagggml_cuda_mul_mat_qhad already computed and already usedto size the
src1padding throughggml_cuda_mmq_get_J_max. The two agreed by construction whilethe only input was
ne01, but any other reason to want the bounds-checked configuration was silentlydropped. The decision is now carried in
mmq_args, so the configuration the padding was sized for isthe configuration that runs.
Measurements
DGX Spark GB10, sm_121a,
nvidia-smi -lgc 300,2100, every cell at 2086 to 2093 MHz, single node withthe GPU lock held, no other compute process resident.
test-backend-ops perf -b CUDA0, each cellbracketed by a
basecell before and after; the two base cells agree to within 2 % on every densecase and to 0.1 % end to end.
Dense
MUL_MAT, TFLOP/s (N x M x K, M = tokens)Batch 32 and q8_0 are unchanged: q8_0, q4_0, q5_0, iq4_xs and iq4_nl are deliberately left on the
stock tile, since the narrow tile was worth under 2 % for them at large N and cost 23 % at N = 512.
The types that gain are the ones whose mainloop is rescale bound (the K-quants) or FP4.
MoE
MUL_MAT_ID, TFLOP/s (256 experts, 8 used)Two base cells are shown because the first MoE cell of a window is cold; the patched cell sits inside
the base bracket on every case, i.e. no regression.
Without the
mul_mat_idguard the same build gives q4_K 1024 x 2048 x 2048 at 13.64 (-18.6 % againsta 16.75 base) and q8_0 at 9.80 (-20.9 %), which is what the guard is there to avoid.
End to end
llama-batched-bench -m Qwen3.8-27B-UD-Q4_K_XL.gguf -c 4608 -b 2048 -ub 2048 -ngl 99 -fa on -npp 512,2048 -ntg 16 -npl 1, base / patched / base in one lock window at one clock state.Things that were tried and did not work
cp.asyncfor the y tile. Both K halves fetched withcp.async.cginto a double buffer, onecp.async.wait_allplus one__syncthreadsreplacing the two register copy loops and one of thetwo barriers per K step. 21 to 24 % slower on every dense shape (q4_K 62.4 -> 48.6 TFLOP/s at
M = 512). The extra y tile in shared memory costs more occupancy than the asynchronous copy buys.
stack on q4_K J = 128); the build faults at runtime on
MUL_MAT_ID.dA_ik * dB_jkis a different scalar for every (i, j, k-block) and no partial sum can stay inint32 across k-blocks. Accumulating in int32 across a 256-wide superblock, which is what
CUTLASS-style block-scaled kernels do, requires quantizing the activations with one scale per 256
values instead of per 32. That changes results and needs a new
quantize_mmqlayout; it is notpart of this PR.
Proof that nothing else is touched
The diff is four files, all under
ggml/src/ggml-cuda/: the Blackwell tile table, the FP4 vec_dot,and the two lines of
mmq.cu/mmq.cuhthat carry the tile choice. No RPC, backend-registry,CPU, Metal, Vulkan, HIP or SYCL file is touched.
Every table entry that changed is inside
ggml_cuda_mmq_get_config_dgx_spark, reached only underan exact
ggml_cuda_highest_compiled_arch(cc) == GGML_CUDA_CC_DGX_SPARKon the host and__CUDA_ARCH__ == GGML_CUDA_CC_DGX_SPARKin device code. Themul_mat_idterm infallbackisgated the same way. Every other GPU selects exactly the configuration it selected before: a
13200-row evaluation of the real host selector (31 compute capabilities x 20 types x 11 J x both
fallbackvalues) shows 0 differing rows against the merge base for every non-GB10 cc, and 0 of13200 on a build whose architecture list does not contain sm_121.
The narrowing took two commits, and this is worth stating precisely rather than as "narrowed".
The first narrowing was entirely in the tile configuration selector. The FP4 accumulator change in
mmq-vec-dot.cuhsits insideggml_cuda_mmq_vec_dot_fp4_fp4_mma, compiled for everyarchitecture defining
BLACKWELL_MMA_AVAILABLE, that is__CUDA_ARCH__in [1200, 1300), which isall of sm_120 too. So sm_120 went from 5 of 23 MMQ translation units differing, to 2, and only
reaches 0 with the second commit that gates that hunk as well. Per-arch, differing translation
units against the merge base, from
cuobjdump -sassover all 23 units at75-real;80-real;90-real;120a-real;121a-real:The severity of that middle column is worth keeping separate from the tile question above: what
sm_120 was getting was an unmeasured code generation change, not an unmeasured numerics
change. The merge base plus that hunk alone, against the merge base, is bit-identical on all 48
(type, shape) cases, MXFP4 and NVFP4 included. Those are different risks. It is gated anyway,
because it was swept on a DGX Spark and nowhere else.
GB10 loses nothing to that second gate: a strict SASS comparison of the tile-selector-narrowing
commit against the current head is byte-identical on sm_121a for all 23 translation units.
The FP4 vec_dot change is inside
ggml_cuda_mmq_vec_dot_fp4_fp4_mma, which only exists for theblock-scaled FP4 MMA path, and is now itself behind
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ == GGML_CUDA_CC_DGX_SPARK, with the merge base'szero-a-temporary-and-add form restored in the
#else.cmake -DGGML_CUDA=OFF -DGGML_RPC=ON -DGGML_NATIVE=ONbuilds clean with the patch applied(
test-backend-opsandllama-benchboth link).test-backend-ops test -b CUDA0on the patched build: 13572/13572 tests passed, Backend CUDA0:OK. Restricted runs:
-o MUL_MAT1193/1193,-o MUL_MAT_ID869/869, both matching the mergebase exactly, single device, no
--rpc.No AMD hardware was available, so the HIP argument is from source only: the device branch is
inside the non-HIP arm of
#ifdef GGML_USE_HIPand additionally testsdefined(__CUDA_ARCH__),and the host branch sits after the unconditional
GGML_CUDA_CC_IS_AMDreturn and is itself gatedon
GGML_CUDA_CC_IS_NVIDIA(cc). The host half is executed evidence: 0 of the 1320 AMD rows in theselector matrix differ.
Still open
The CUTLASS W4A4 path of ggml-org#26704 reaches roughly 250 TFLOP/s on this class of GPU,
still about 1.75x what this gets to, so the tile table is not the end of the story. ggml-org#26977
(routed-MoE tile width from the mean columns per expert) and ggml-org#26159 (compact expert/tile job
list for the NVFP4 MoE grid) cherry-pick cleanly onto this and are a real MoE win here (+37 % q4_K
experts at 128 tokens, +39 % NVFP4 at 2048); they are upstream PRs, so they are not duplicated in
this one.