The NVFP4 TMA route reads activation scales as one tile: operator 0.81x to 0.87x, prefill +4.1%, bitwise identical - #160
Conversation
The activation scale descriptor is a 16-byte box over the scale plane, so one stage of the TMA pipeline asks for BlockM separate 16-byte transfers to deliver 4 KiB. The weight codes beside it travel as one 64-byte box per row and cost nothing like that. A standalone harness that runs the producer alone - the same four descriptors, the same 29,696-byte stage, the same three-deep mbarrier ring, no math - shows where the stage time goes: the codes alone move at 6977 GB/s; adding the weight scales costs 14.8% more time for 4.2% more bytes; adding the activation scales costs 73.0% more time for 16.7% more bytes. Requesting the same bytes as one [BlockM tokens, 16 groups] tile restores the codes-only rate, 7009 GB/s against 4732 for the same payload. The byte order inside a tile is unchanged, so the shared image the consumer reads is identical and the consumer is untouched: only the address the quantizer writes to, the descriptor shape and the request coordinate move. The quantizer writes the plane blocked when the GEMM that follows will take the TMA route and keeps it row-major otherwise, because the small-T and GEMV routes read the same buffer. The predicate is the route's own; the fused SwiGLU route admits every multiple of 256 from 256 up, which is wider than the shared one. RTX 5090, sm_120a, 510 W cap, 3105 MHz boost, CUDA 13.1.115, Release. Both builds come from one tree, interleaved ABAB, three rounds, one process per run. Operator, T=4096, NVFP4 A4: shape before us after us ratio GB/s TFLOP/s linear_add [5120, 17408] 759.808 612.352 0.806 364 -> 452 961 -> 1192 linear_add [5120, 6144] 303.104 262.144 0.865 501 -> 580 850 -> 983 linear_swiglu [34816, 5120] 1456.128 1214.464 0.834 196 -> 234 1003 -> 1202 Against the hardware limits: the GDDR7 spec peak is 1792 GB/s, so these shapes sit between 11% and 28% of it before and between 13% and 32% after - the route is nowhere near bandwidth-bound at either end, and what the change buys is transaction width rather than bandwidth. For the tensor core, the block-scaled NVFP4 MMA measures 2117 TFLOP/s in isolation on this part, and 1897 when fed in this kernel's own ratio of 4 ldmatrix.x4 and 8 ldmatrix.x2 per 32 MMA; against that 1897 the three shapes move from 51%, 45% and 53% to 63%, 52% and 63%. Those two ceilings were taken at full boost and these runs did not sustain it, so the tensor-core fractions are a lower bound; the bandwidth fractions are not affected, memory speed being independent of the SM clock. Controls in the same runs, below the 256-token floor where the route is not taken: linear_add at T=64 and T=128, fused SwiGLU at T=64. All three are 1.000 to the printed digit, which is what shows the two builds are otherwise identical. End-to-end confirmation, Qwen3.6-27B NVFP4, 32768-token prompt, prefill chunk 4096, KV int8: prefill 9426.4 -> 9811.8 tok/s (+4.09%), 3.4762 -> 3.3397 s, with round-to-round spreads of 101.4 and 98.6 tok/s. Decode is the untouched control at 87.945 -> 87.958 tok/s (+0.01%): its token counts never reach the tile. Resource deltas. Workspace capacity and peak, weight capacity and KV payload are byte-identical between the builds as the engine reports them: 639,926,272, 17,206,931,200 and 1,107,296,256. No kernel changes registers or shared memory, and none of the 37 functions in the touched translation units spills. The binary gains three kernel instantiations: the quantizer template takes a flat/blocked flag, so each of the three activation geometries is emitted twice, at 40 registers and no shared memory. Numerical validation. Values do not change - only the address a scale byte is written to - and greedy generation matches the previous binary byte for byte on four prompts of 8K, 64K and two of 55K tokens. ctest is 102 of 102, with six cases skipped for want of real model artifacts. clang-format 23.1.0 reports eight violations across these files. All of them are on lines this change does not touch, so they are left as they are.
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. |
|
what about T=1024? |
check the RTX5090 spec for the number, I remeber it's 1736 TFLOPS for nvfp4 |
AH 3105 MHz boost ok |
|
T=1024 is where the route turns on and it is the shipped default, so it should have led the table. Same tree, interleaved, three rounds, ratio new/old: T=768 is the control below the floor. End-to-end at chunk 1024, same 32768-token prompt: 8607.6 -> 9123.5 tok/s (+5.99%, spreads 110 and 95) against the +4.09% I reported at 4096; decode unchanged. T=1280 is the sweep's weakest point - five 256-token tiles sit on the grid worse than their neighbours. Happy to fold these into the PR body if you would rather have them there than in the thread. |
…s (from MichaelDementii/ninfer PR Neroued#160 545f64b)
…pair Neroued#167/Neroued#160, dylan wave 2 (7dd98fd), gzenz safety net (08636ed), md fp8-KV campaign, T22-T27 plan
|
A correction to one number in the body, before it trips anyone up. I re-measured the NVFP4 tier directly: back-to-back MMAs with operands in registers, the result On the 5090 I have in front of me (3105 MHz nominal boost, 525 W cap) the highest SM clock I can The three percentages in the table divide by the fed ceiling of 1897, not by 2117. That one was |
…(port PR Neroued#160) Cherry-pick 545f64b (MichaelDementii): activation-scale TMA descriptor was BlockM separate 16-byte requests per 4KiB tile; now one [BlockM tokens, 16 groups] tile. Byte order within a tile unchanged -> consumer identical; quantizer writes blocked only when the following GEMM takes the TMA route (small-T/GEMV stay row-major), BlockM=256 everywhere (verified on our 27B geometry: TmaM256N128). Measured by author on Qwen3.6-27B NVFP4: prefill +4.09% (32k prompt, chunk 4096), decode +0.01%, greedy bit-identical, ctest 102/102. Our A/B: pre-A 178.5 tok/s / 10.2k-3.5k prefill / needle 12/12 (2026-09-07).
|
Independent verification on our rig (cometkim Qwen3.8-27B nvfp4full, our fork at upstream 487f897, commit 7e8ad2e): Same probes, same session, before/after (cherry-pick applied cleanly, zero conflicts on 487f897):
Context for the smaller-than-your-claims deltas: our server runs Also note: Happy to share the config ( @Neroued — this PR is safe to merge: bitwise-identical greedy output on our 8K/55K/64K prompts + ctest 102/102 (as you stated), and it improves prefill at the shipped default chunk 2048 too. |
… tooling') main is upstream 487f897 + 9 commits: own nvfp4full registration + graft tool, the Neroued#148/Neroued#97/Neroued#160/Neroued#61 PR ports (temp, drop as they merge), the registry doc, and curl for healthchecks. Registry table in PATCHES.md now covers all 9 (added curl + registry rows, date -> 2026-09-15).
3 participants
The activation scale descriptor is a 16-byte box over the scale plane, so one stage of the TMA
pipeline asks for BlockM separate 16-byte transfers to deliver 4 KiB. The weight codes beside it
travel as one 64-byte box per row and cost nothing like that.
A standalone harness that runs the producer alone - the same four descriptors, the same 29,696-byte
stage, the same three-deep mbarrier ring, no math - shows where the stage time goes: the codes alone
move at 6977 GB/s; adding the weight scales costs 14.8% more time for 4.2% more bytes; adding the
activation scales costs 73.0% more time for 16.7% more bytes. Requesting the same bytes as one
[BlockM tokens, 16 groups] tile restores the codes-only rate, 7009 GB/s against 4732 for the same
payload.
The byte order inside a tile is unchanged, so the shared image the consumer reads is identical and
the consumer is untouched: only the address the quantizer writes to, the descriptor shape and the
request coordinate move. The quantizer writes the plane blocked when the GEMM that follows will
take the TMA route and keeps it row-major otherwise, because the small-T and GEMV routes read the
same buffer. The predicate is the route's own; the fused SwiGLU route admits every multiple of 256
from 256 up, which is wider than the shared one.
RTX 5090, sm_120a, 510 W cap, 3105 MHz boost, CUDA 13.1.115, Release. Both builds come from one
tree, interleaved ABAB, three rounds, one process per run.
Operator, T=4096, NVFP4 A4:
Against the hardware limits: the GDDR7 spec peak is 1792 GB/s, so these shapes sit between 11% and 28% of it
before and between 13% and 32% after - the route is nowhere near bandwidth-bound at either end, and
what the change buys is transaction width rather than bandwidth. For the tensor core, the
block-scaled NVFP4 MMA measures 2117 TFLOP/s in isolation on this part, and 1897 when fed in this
kernel's own ratio of 4 ldmatrix.x4 and 8 ldmatrix.x2 per 32 MMA; against that 1897 the three shapes
move from 51%, 45% and 53% to 63%, 52% and 63%. Those two ceilings were taken at full boost and
these runs did not sustain it, so the tensor-core fractions are a lower bound; the bandwidth
fractions are not affected, memory speed being independent of the SM clock.
Controls in the same runs, below the 256-token floor where the route is not taken: linear_add at
T=64 and T=128, fused SwiGLU at T=64. All three are 1.000 to the printed digit, which is what shows
the two builds are otherwise identical.
End-to-end confirmation, Qwen3.6-27B NVFP4, 32768-token prompt, prefill chunk 4096, KV int8:
prefill 9426.4 -> 9811.8 tok/s (+4.09%), 3.4762 -> 3.3397 s, with round-to-round spreads of 101.4
and 98.6 tok/s. Decode is the untouched control at 87.945 -> 87.958 tok/s (+0.01%): its token counts
never reach the tile.
Resource deltas. Workspace capacity and peak, weight capacity and KV payload are byte-identical
between the builds as the engine reports them: 639,926,272, 17,206,931,200 and 1,107,296,256. No
kernel changes registers or shared memory, and none of the 37 functions in the touched translation
units spills. The binary gains three kernel instantiations: the quantizer template takes a
flat/blocked flag, so each of the three activation geometries is emitted twice, at 40 registers and
no shared memory.
Numerical validation. Values do not change - only the address a scale byte is written to - and
greedy generation matches the previous binary byte for byte on four prompts of 8K, 64K and two of
55K tokens. ctest is 102 of 102, with six cases skipped for want of real model artifacts.
clang-format 23.1.0 reports eight violations across these files. All of them are on lines this
change does not touch, so they are left as they are.