Repository navigation
glm cuda: the L2 prefetch's ffn-head site also reads a MoE layer's shared gate-up, after its router - #3156
glm cuda: the L2 prefetch's ffn-head site also reads a MoE layer's shared gate-up, after its router#3156Bizuayeu wants to merge 1 commit into
Conversation
…ared gate-up, after its router Site f read a MoE layer's router alone. The shared expert's gate-up runs on the same window once the routed experts are done, and its scales and codes now follow the router in the site's rows, inside the same budget. The table's test gives layer 1 a block FP8 shared gate-up and checks where the 12 MiB budget cuts it.
bbea1fe to
1dd4d08
Compare
|
Thanks for the fresh head and the careful receipt. Adding the shared gate-up to site f is a small, clear change. Reading its scales and codes after the router, inside the same budget, is easy to follow, and the table test now checks both the added rows and the cut when space is tight. The new head applies cleanly. Format, diff check and lean are clean, with zero new lean problems, and the changed file stays well under the size limit. The prefetch computes nothing, so served tokens are unchanged by construction. Verdict: DEVICE. The speed gain and token hashes need a device check, and we will run that on our side. The changelog line can go under Unreleased when this lands. Must fixes: none.
|
|
Our device run on two GB10s at TP=2, main at the merge base against this head (1dd4d08), GLM-5.3-Flash Jammed XSFP4, fresh launches base, new, new, base: GREEN.
Thank you @Bizuayeu. It's going into the landing queue.
|
On two GB10s at TP=2, Jammed XSFP4, one stream at a 40,960-token window, greedy,
mainagainst this branch, the mean of three launches each, run alternately:mainmainttft_s)ttft_s)Decode is tokens per second, so higher is faster; prefill is seconds to the first token, so lower is faster. Decode is about 1% faster, each prompt beyond the spread of either side's three launches, with the same tokens and no new switch; prompt processing stays within that spread, as the prefetch does not run there.
What this changes
TF_GLM_L2PF's site f, before
gather_aat a layer's ffn head, read a MoE layer's router alone. It now also reads the layer's shared expert gate-up after the router, its scales and then its codes, inside the sameTF_GLM_L2PF_MBbudget. When the prefetch was planned in #1580, site f left the shared expert out, on the guess that the routed experts' reads would evict it from L2 before it runs; at TP=2 on GB10, the table above says reading it anyway is faster. At TP=2 a rank's block FP8 shared gate-up is 8 MiB and its router 2.25 MiB, so both fit the default 12 MiB; with less room the budget cuts the gate-up's codes, as the table's test checks. On GLM-5.3-Flash at TP=2 the start-up line readsTF_GLM_L2PF: 310 rows, 1423 MiB a decode windowagainstmain's 226 rows and 1066 MiB. Replies are identical.mainat223da7659b, one commit: the change (1dd4d08b73). Its CHANGELOG line waits for landing:mainreleased 1.2.8 after this base and the line no longer merged into its CHANGELOG, so it can go under Unreleased when this lands.atc overlap --json:complete: false(pullFilesUnavailableand its time and count limits), with no overlap for this branch (worktreeOverlapsis empty). Throughgh pr list --json files, none of the 359 open pull requests changescuda_l2pf.zig.Checks on this head
1dd4d08b73. Every receipt is from its binaries.python3 tools/zig/lean_check.py, on the base and on this head: 341 problems on both, 0 new.tools/pr_check.py --base origin/main --pr 3156: 0 errors and 1 warning, that no Unreleased line names this pull request, as the line waits for landing; lean 341 against 341 with 0 new.zig build test-cuda-host, 6,392 of 6,422 with 30 skipped (393 of 393 steps), withzig fmt --checkandzig ast-checkon the changed file. With the test's change alone onmain,test-cuda-hostfails the table's test and nothing else (6,391 of 6,422, 30 skipped, 1 failed). CI's pytest from the root,-q -rs: 450 passed, 86 skipped, 519 subtests passed;main's CI at1e953d8400gives 452 passed and 84 skipped, and the two more skipped here aretools/zig/test_nfkd_tables.py's, under Python 3.12's Unicode 15.0.0. On two GB10s with this branch'stf-cuda-teston Jammed XSFP4 at TP=2:glm-lanes-cover46 PASS,glm-lanes-concurrent38 PASS andglm-shared-check1 PASS on rank 0, with rank 1's 1 PASS each and both ranks exiting 0; their logs show the prefetch on with this branch's 310 rows.1dd4d08b73overmainat09a9eb3317: fmt, zig, zig-windows, zig-macos-dev-tools, pytest, benchmark-clients on Ubuntu and Windows, and release-paths; host-nightly skipped, as on every pull request.maindoes not read yet. Other CUDA chips. With nvfp4 experts cuda: decode up to 24% faster (GLM-5.3 83.8 -> 104.1 tok/s), same bits, via a per-warp cp.async ring #2266's decode kernels, not merged, which change what the routed reads leave in L2.Receipt
main223da7659band this branch's1dd4d08b73(the head), each built with an empty Zig global cache, Zig 0.17.0 and nvcc 13.3 for sm_121. Two GB10s over a direct link (driver 580.178.04),Bizuayeu/GLM-5.3-Flash-Jammed-XSFP4@cae94e69from a local copy (each of its 132 safetensors files has the revision's sha256; the local index differs in its metadata alone), a 40,960-token window, one stream, vision off.cuda_l2pf.zig's sources for site f on MoE layers, read on every decode window's side stream. Ran at TP=2 on GB10; the rest is under "Not run".tools/bench_concurrent.py --alone --serialortoken_shaagainst"draft": false: the three decode-check prompts, 512 greedy tokens each, give the sametoken_sha256drafted and with"draft": false, onmainand on this branch:ae95484b0c71,315eaea6c98band461c6e99e77f. Every request's output sha in rank 0's done lines is the same, in order, on both.tools/bench_openai.py: as of0d3ad792f, its two prompts with"draft": falsein each request, 256 tokens, five repetitions, one stream at a 40,960-token window, each launch's median, the mean of three launches: the table. This branch decodes 0.98% and 1.24% faster on the two prompts, and faster in each of the three alternated pairs (+1.38%, +1.13% and +0.45%; +1.31%, +1.14% and +1.27%), while each side's three launches spread by under 0.9%. Both sides' load timing picked the same expert form.ttft_sfromtools/bench_cache.py --require-coldas of19d1a9af6, on the decode-check prompts' tokens repeated to 9,600 and 38,400, the mean of three launches: the table. Prompt processing is within 0.3% ofmain(+0.25% and +0.22%), inside the spread of each side's launches (0.5% and 0.6%); the prefetch does not run there.Checklist
1dd4d08b73, one commit onmainat223da7659b: every receipt is from its binaries.)main, fails and no other test does; it needs no GPU.)glm-lanes-cover,concurrentandsharedpass at TP=2 with it on.)"draft": false, a resumed prompt equals a fresh one, and concurrent equals solo. (Drafted equals plain on both; resumed and concurrent by the lanes tests.)--drafter, no slower than without it, greedy and sampled. (One stream, greedy, MTP drafts off, on Jammed XSFP4: decode 1.0% and 1.2% faster, prompts within 0.3% ofmain, means of three launches each. The widest admitted and--drafterwere not measured.)