Repository navigation
Conversation
…byte load The fact entries loaded rows r and r + 8 as two 8-byte pairs, and the quad (a0..a3) = (xa.x, xb.x, xa.y, xb.y) was rebuilt with MOVs before every mma: 644 MOVs in the 4-tile k-step of tf_nvfp4_fact_swiglu (2,291 instructions, 17.9 per HMMA). On items with all 4 row tiles (WM 1, so the CTA's warps share the layout), X is now staged with the two rows interleaved word by word: pair p, unit u holds one quad, stored at u xor 5 on odd pairs so both the quad reads and the stores are conflict-free. The copy goes through registers (ld.global.nc 16 bytes from each row, two 16-byte stores), loaded a whole k-step ahead and stored after the step's math. The arithmetic is unchanged: same values in the same k-slots, same mma from zero, same fma order. Items with fewer tiles keep the cp.async copy and the pair reads; the prompt_* entries are untouched. sm_121a, nvcc 13.0: the 4-tile k-step drops to 1,772 instructions (13.8 per HMMA), MOV 119. SASS of the other 55 nvfp4_experts functions equals main's on sm_80/86/89/120/121 (13.0 and 13.2). Records re-derived with kernel-contract-limits for 13.0 (registry.cuda) and 13.2 (a cuda_132 override for the fp32 down); the 13.3 override of the down, measured before this change, is dropped until 13.3 re-derives it. tf-cuda-test nvfp4-fact passes (bytes equal to the decode form at 1 to 300 rows). Synthetic Flash Next layer (D 2560, NI 640, E 512, top-10), 4 cold rounds: fact_swiglu at T 8192 9.87 -> 9.05 ms, fact_down_f32 -1.3 % (Nemotron shape -3.3 %), T 512 and 1536 flat, T 2048 fact_swiglu +0.9 %; output bytes equal in every round. Flash Next FP4 served at the default 1,536-row chunk: first-token logits equal at 8k and 32k, cold TTFT 8k 4.12 / 4.18 s and 32k 16.60 / 16.53 s (main / this, 2 rounds), texts equal.
…back on today's code On an RTX PRO 4500 (sm_120a) the down entry measured a reproducible +1.8 % at T 2048 with balanced routing (NT 3 items, 141 -> 192 registers), so by ashhart#1479's rule the down keeps today's code and only fact_swiglu changes: XQ now also requires M == 2. SASS (nvcc 13.0 and 13.2, sm_80/86/89/120/121), nvfp4_experts against b30d5b9: only tf_nvfp4_fact_swiglu differs, on all five; the other 60 functions, tf_nvfp4_fact_down_f32 among them, are identical. Against the previous head ea5686d only tf_nvfp4_fact_down_f32 differs, so fact_swiglu is the code measured there. The down's records are main's again (13.0 row, no 13.2 override, the 13.3 override restored).
…r more cubins The timer ashhart#1479's synthetic tables came from: tf_nvfp4_fact_swiglu and tf_nvfp4_fact_down_f32 on a Flash Next expert layer (D 2560, intermediate 640, 512 experts, top-10, random NVFP4 weights, uniform routing, items of up to 64 pairs), each kernel captured in a CUDA graph and replayed, modules in alternating order. Every module's output bytes are checked against the first one's; the exit status is 1 if any differ. Needs torch with CUDA; the driver is called through ctypes.
GB10 served check: Flash Next FP4-Local, main vs #1525Main Per chunk: 4 server boots in ABBA order, 2 per arm. Each boot had an untimed warmup at every length, then 3 timed requests per length. That is 6 timed requests per arm per cell, but only 2 independent boots. Another job paused the run once between boots 1 and 2. The interrupted boot was discarded whole and the run resumed with the same binaries and settings. Correctness
Median TTFT, main → PR
Reading
Reproducibility bundle (harnesses, fixtures, per-boot summaries, all timing observations, 1 Hz GPU telemetry, analyzers, source/model/binary hashes, rebuild steps): https://github.com/Maxrio9/TensorFold/releases/download/benchmark-1479-sm120-20261028/tensorfold-1525-gb10-served-results.zip |
|
I'll follow up with a revised kernel version to eliminate the slow cells. We also need to rule out thermal throttling during benchmarking by introducing cooldown periods between tests. Thermal throttling can significantly affect performance and undermine the reproducibility of benchmark results. |
|
Thanks @jschmied. The source keeps every MMA starting from zero, the K-group and half order, the scale-FMA order and the epilogue. A symbolic check of the new load layout passes for item counts 49 to 64 and odd K tails. The open question is speed at the default chunk. Your newer served receipt has TTFT up 0.64% at 8k and 0.52% at 32k on the default chunk, with both boot pairs agreeing and the 8k ranges not overlapping, while the 8,192-row chunk gains about 2%. Flash Next takes Once the default holds, we'll run tokens and first-logit bytes, alternating served boots at both chunk sizes, the five-SM limits, and per-function SASS identity for everything except |
|
Closing this. The new layout doesn't pay at the default chunk, and I found out why. Cause: weight reuse between neighbouring items, not the X path itself. I dumped the served plans (items, counts, members) of one 8,361-token prompt at the default chunk and replayed them in the timer, with the same
Each of these changes alone leaves the sign the same:
Changing only the expert ids flips it:
So the new layout wins when several items of one expert run side by side and read its weights from L2, and loses when each item streams its weights from DRAM. Under real routing at the 1,536-row chunk, an expert with four-tile items has about 1.5 of them, so the default path loses; at the 8,192-row chunk most experts carry several items, hence the gain there. I haven't measured the stall mechanism itself. Most likely the warp's wait on the X loads hides less of the weight latency. A cutover keyed on items per expert would only help non-default chunks, so it isn't worth a second code path. Keeping weights resident across one expert's items would help the default chunk too; that would be a separate proposal if the numbers hold up. |
Part of #1479.
What
tf_nvfp4_fact_swigluis the prompt expert kernel Flash Next's expert tile runs for every prompt pass of 64+ rows. On main it runs 17.9 instructions per tensor-core mma (sm_121a, nvcc 13.0), and 644 of the 2,291 instructions in its 4-tile k-step are MOVs. Rows r and r + 8 come in as two 8-byte loads, and the MOVs rebuild the mma's A quad from them before every mma.On items with all 4 row tiles, the branch stages X with the two rows interleaved word by word, so each quad is one 16-byte shared load. Pair p, unit u holds one quad, stored at u xor 5 on odd pairs, so loads and stores are conflict-free. The copy goes through registers, loaded a whole k-step ahead and stored after the step's math. Items with fewer tiles keep the cp.async copy and the pair reads. The path requires
M == 2, sotf_nvfp4_fact_down_f32and theprompt_*entries keep today's code.fact_swiglustays at 255 registers.Commits:
ea5686db6(the quad path for both fact entries),98b08b2d5(back to today's code forfact_down_f32after the RTX PRO 4500's +1.8 % at T 2048 balanced, per the rule in #1479), and1e627ed14(the timer undertools/zig/). Baseb30d5b9d3. Not rebased, because it applies to current main as is.Callers. At this base the only production caller of
runFact(which launchestf_nvfp4_fact_*) is Flash Next's expert tile,zig/src/families/flashnext/cuda_nvfp4_tile.zig. Bolt's IQ2_XXS experts run on the ggml kernels, and NVFP4 experts on the grouped prompt path go throughrunPrompt(tf_nvfp4_prompt_*). Flash Next is the only served model this reaches.SASS
nvfp4_experts, per function, sm_80/86/89/120/121, nvcc 13.0.88 and 13.2. Onlytf_nvfp4_fact_swigludiffers, on all five SMs. The other 60 functions (12 per SM,tf_nvfp4_fact_down_f32among them) are identical. This holds against the branch baseb30d5b9d3and against current maind0c9dc9fb, merged with the branch. No other image includes a changed file.Command:
tools/zig/cuda_sass_identity.pyfrom #1440 (not on main yet). It builds every catalog image a change reaches from both trees, with each tree's catalog flags, and comparescuobjdump -sassper function:Without that tool: build
nvfp4_experts.cuper tree withnvcc -cubinfor each SM, then diffcuobjdump -sass -fun <symbol>per function.Records
zig/src/core/contract/nvfp4_fact_records.zig:fact_swigluis 255 registers on all five SMs (main: 249 on sm_80/86).fact_down_f32is main's again: one row for nvcc 13.0, no 13.2 override, and the 13.3 override restored innvcc_133_records.zig.kernel-contract-limitson all-arch fatbins confirms both fact records as current on all five SMs under nvcc 13.0. Under 13.2fact_swigluis current.fact_down_f32reads 142 registers on sm_120/121 against main's 141 and is reported stale (it still passes). That is main's record and main's code, so it is unchanged here.Timer
tools/zig/nvfp4_fact_timer.py: the synthetic-layer timer the tables below came from, as posted in #1479. It runs Flash Next's routed-expert layer (D 2560, intermediate 640, 512 experts, top-10, random NVFP4 weights, uniform routing, items of up to 64 pairs). Each kernel from two or more cubins is captured in a CUDA graph and replayed, with the order alternating by round. Every cubin's output bytes are checked against the first's, and the exit status is 1 on any difference. Build and run commands are in--help.Measurements
fact_swiglu, synthetic layer, base → branch. Positive means slower.The GB10 column was measured on
ea5686db6, whosefact_swigluSASS is identical to this branch's. The RTX PRO 4500 column is @Maxrio9's measurement ofea5686db6againstb30d5b9d3in #1479.On GB10 the kernel is bound by its weight stream at Flash Next's default chunk. A diagnostic build with zero copy cost (wrong numbers, timing only) equals base at 1,536–2,048 rows. Served Flash Next (FP4-Local, default chunk, GB10, on
ea5686db6) was level: cold TTFT 8k 4.12 / 4.18 s, 32k 16.60 / 16.53 s, first-token logits byte-equal. The gain should show atTF_FLASHNEXT_CHUNK=8192, and possibly at the default chunk on a card with more bandwidth per SM.Checks (head
1e627ed14, GB10, nvcc 13.0)tf-cuda-test nvfp4-fact: pass. Gate-up and fp32 down are bit-equal to the decode form at 1/7/17/22/64/65/130/300 rows, FP64-bounded, and the control goes red.kernel-contract-limitson all-arch fatbins (-Dsm=80,86,89,120,121), nvcc 13.0 and 13.2: both fact records pass on all five SMs (13.2's down note is under Records). The only failures are 10qwen35 tf_q35_attn_partial(8)1024-thread launch records, in files this branch doesn't change.zig fmt --check zig: clean.python3 tools/zig/lean_check.py: 0 new (428 on the base, 428 on the head).zig build test -Dnvcc=/usr/local/cuda/bin/nvcc: RC 0, 0 errors.git merge-treeshows it applies cleanly to maind0c9dc9fb. The only open PR that touchesnvfp4_experts.cuandnvfp4_fact_records.zigis my nvfp4 experts: Flash Next decode rows in equal waves, two warps a unit for one row #1351 (Flash Next decode rows), and it also merges cleanly with this branch.Proposed check
As in #1479: the kernel contract, plus served Flash Next FP4-Local against main on a GB10. Cold TTFT at 2k, 8k and 32k, at the default chunk and at
TF_FLASHNEXT_CHUNK=8192, with replies and first-token logits byte-equal.