Repository navigation
Conversation
|
Updated the description: the Speed section now gives the comparison with the Triton set run on this head ( |
|
Hello @Bizuayeu, Thank you for the GLM-5.3-Flash CUDA family. On this head it calls Your body reports a serve on two GB10s at Tensor parallel over hosts is what the code does. You report the server at TP=3 was not run, because a ring of /32 addresses gives the other ranks no bootstrap interface, and that the engine checks at TP=3 passed on 1.0.8. The comment at The body lists #805 among the commits under review. Four things have to change before this is ready.
Lines past 120 columns, the longest at Once those four are done, the gate is this family serving the NVFP4 checkpoint on a CUDA GPU. More than one GPU needs a second card and NCCL. The receipts are a drafted reply equal to Thank you again. |
…ecodeF32 and promptF32) A tensor-parallel family sums its ranks' partial rows before rounding, so it needs a projection's fp32 sums as the kernel holds them. qmmf.cu and prompt16.cu already take F32 as a template argument; this instantiates FP4's fp32 tiles (and reduce_kernel<true> for the split path), loads them beside the bf16 ones, and adds Lane.matmulOut, Prompt.matmulOut and qlinear's decodeF32 and promptF32. Other formats refuse fp32 out with EpilogueNotBuilt, as the experts do. matmul, decode and prompt keep their signatures and bf16 bits.
…n block The GLM-5.3-Flash CUDA port reads nvidia/GLM-5.3-Flash-NVFP4 as stored. admit() checks config.json's quantization_config (top level, else text_config's) for quant_method modelopt, quant_algo NVFP4 and 4-bit float weights in groups of 16, and refuses anything else with a reason the caller logs (the Why of nemotron/config.zig), so a test reads the reason instead of failing on an error log. MIXED_PRECISION is refused for now. Which linears are NVFP4 is left to the tensors, not to the block's ignore list: the pinned revision 423acf37 stores the MTP layer's experts in BF16 without listing layer 45 there (da920bb0 added it). The fixture is that revision's config.json.
nvfp4.zig is the format the GLM-5.3-Flash CUDA port's W4A16 linears read: e2m1 codes two to a byte (input 2j in byte j's low nibble) [n, k/2], an e4m3 scale per 16 inputs [n, k/16], and weight_scale_2 as an fp32 scalar that multiplies. dims() checks the four tensors fit together; Linear.read() maps them from a Checkpoint and marks input_scale as used, since the W4A4 math's static input scale is not read by W4A16 but a leftover tensor would refuse the checkpoint; row() gives a row as e2m1 x e4m3 x global in f32, the reference the GPU kernels will be compared with. stored() says a module is NVFP4 by its weight_scale tensor. The e2m1 and e4m3 values come from core.quant's nvfp4 and fp8 (ashhart#730), as f32. The fixture is the whole header of nvidia/GLM-5.3-Flash-NVFP4 (147,661 names, dtypes and shapes, gzip'd): every NVFP4 linear reads, 42 MoE layers x 288 experts x 3 plus the three dense MLP layers' 3, all under .mlp., and the MTP layer's experts stay BF16.
…e with them A GLM-5.3-Flash rank reads about half of the 183 GB checkpoint through the mapping, and its column cuts touch whole tensors. On GB10's unified memory the page cache holds memory the rank's device allocations need (our Python serving engine drops it with posix_fadvise for the same reason). safetensors.File.release drops a file's pages from the mapping (madvise) and the page cache (fadvise DONTNEED); Checkpoint.release does every file. The mapping stays valid and refaults on a later read, which a test checks. The GLM-5.3-Flash CUDA family (ashhart#548) calls it after each layer it loads.
…me from its own declarations The CUDA host wired Nemotron's prompt-cache copies, rules and budget in by name. They now come from the family's `snap_*` declarations through `cuda_snaps.zig`: a family without `snap_save` gets no prompt cache, `snap_ready` sets aside its state buffers, and `snap_grid` makes the store keep states only on the family's chunk starts (a family whose rows depend on where a prompt's chunks begin), and `warm_quiet` turns on the quiet-turn reply prefill (warm turns and the store's `.waits`). Nemotron declares `warm_quiet` and no grid and keeps 1.0.9's rules: any row, the head one token on, warm turns, the same budget. Three more hooks for a family split over hosts: `weightShare` (the weights this rank holds, for the memory check), `explains_open` (the family's own text for an open error) and the learner only for a family with `slide`. Nemotron's open, cache, warm turns and --slide are unchanged; `cuda_snaps.zig`'s test covers the rules each way.
…skip), as the decode form does A family that computes one expert itself (GLM-5.3-Flash stacks its BF16 shared expert as the last expert id) plans every pair together and has the grouped kernels skip that id. The decode form already takes skip; Grouped.prompt refused it with OptionNotBuilt for NVFP4. The prompt entries now take skip and return before staging any rows of that expert's items. Nemotron passes -1, so its bits do not change. quant-ref now checks it: with the shared expert skipped, its pairs' prompt rows keep what was there and every other pair's row equals the run without skip, for fp8 and nvfp4.
packGateUp packed each matrix into a scratch half and copied the halves into their interleaved places, two allocations and a copy per expert. A loader packs every routed expert of every layer this way (GLM at TP=2: 45 x 256 experts, 2.4 MB a half), and the fresh pages and the copies put about 20 s on the load. packOne's loop now takes the matrix's place among the stacked parts, so packGateUp writes gate and up straight where they go, and it needs no allocator. The words are the same: a host test checks packGateUp against packOne's blocks interleaved.
… over hosts The family ashhart#593 tracks, on the shared NVFP4 layer (`qlinear`, `experts.Grouped`) and hand-written `.cu` for the rest: KDA and MLA/DSA attention, the hyper-connection glue, MoE routing, the MTP draft head and sampling. Rank 0 serves through `tensorfold-native serve` and the other ranks follow its steps; each rank holds its share of every axis (uneven where a width doesn't divide). A prompt piece with at least as many rows as ranks is split over them (`TF_GLM_PREFILL_REDUCE=gather` all-gathers it instead): each rank sends the others their rows' partials and receives its own over NCCL send and receive (`nccl.zig` binds `ncclSend` and `ncclRecv`), and the hyper-connection step adds them in rank order. Drafts and concurrency go through the lanes core: `prefill_step` takes a prompt a chunk at a time between the other streams' rounds, and `hidden_rows` lets several streams' verify windows share one forward. The prompt cache keeps states on the family's 2,048-row chunk starts (`snap_grid`). Checks in `tf-cuda-test`: `glm-ref` holds every projection whose reduction order is the kernels' own against float64 on synthetic inputs; the lanes checks (cover, concurrent, shared) run at TP over hosts and, with `glm-synth`'s two-layer checkpoint, on one GPU.
… EXL3, 64 otherwise) An EXL3 matrix splits only at 128 columns, so at TP=4 the head's 154,880 rows go 303, 303, 302 and 302 blocks of 128 instead of 605 units of 64 each. `Config.head_unit` is 64 by default and 128 when `quantization_config.quant_method` is `exl3`; `Share.of` cuts the vocabulary in it. NVFP4 keeps 64, so its shares at TP=2 and TP=3 are unchanged. The patch BHCC2025 posted on ashhart#593, with its test of the TP=4 and TP=2 shares.
…e width (TP=3's last rank, 640) At TP=3 the routed and shared experts' 2,048 rows split 704, 704 and 640. The block was gcd(512, w), as glue.py picks it, and only 512 and 64 had kernels, so the last rank's 640 (gcd 128) stopped the forward with NoSwigluBlock. The values are element-wise and each 64-column sum is one group whichever block runs it, so the 64-column kernel serves any width that 512 doesn't divide; TP=2's widths keep the 512-column kernel.
…it reads the config The family read config.json with config.parse alone, so a quant_algo other than NVFP4 was never refused on the serve path. Engine.init now runs modelopt.admit on the same text first: a refusal logs admit's reason with the file's path and returns its error (NotModelOpt or UnsupportedQuant), which explain words for the server. glm-synth's two-layer checkpoint writes NVFP4 linears and now says so in its config.json, with the block nvidia/GLM-5.3-Flash-NVFP4 carries, so the lanes checks on it are admitted the same way.
…tes rank 0 admits The server admitted GLM-5.3-Flash's streams and its prompt cache by rank 0's room alone: a follower never returns from open, so a rank with less free memory than rank 0 (another job, the page cache) could run short after the server said it was up. Now the native host passes the family a reading of its budget (Options.room: free memory less the reserve, under TENSORFOLD_CUDA_MEMORY_LIMIT_GB, past what the load took), every rank reads it once its weights are in, and the ranks all-gather the values before a follower starts following (cuda_ranks.gather, which agree now uses too). Rank 0 returns the smallest of the others (Loaded.peer_room) and the host admits by it when it is smaller than its own, naming that rank in a refusal. Rank 0's per-stream and per-state bytes are the largest of the ranks, since the cut gives a remainder to the first ranks, so the room it admits by holds on every rank. Before the load, a rank's weights were taken as the checkpoint over the world size; TP=3's uneven cut puts more than that on rank 0 (22 of 64 KDA heads, 704 of 2,048 expert rows). weightShare now reads config.json and takes the checkpoint's bytes times the rank's largest share of an axis (Share.part), an upper bound but for small replicated tensors. Nemotron declares neither field and takes the host's path as before.
53309a7 to
6380c8d
Compare
|
@ashhart, thank you for the review. The branch is now on 1.0.10 at Speed tables: every row is from On the follower: every rank is its own
Long lines: I wrapped the CUDA lines past 120 columns in the files the branch adds, The description changes with this push:
|
|
@Bizuayeu, thank you for the rebase onto 1.0.10 and for closing the follower-room gap: rank 0 admitting by the smallest room and naming the rank that refuses is just right. A note on where this sits. More GLM-on-CUDA work opened today: @justincasey's #838 and #839 (EXL3 at TP=2, with real speed work) and @BHCC2025's EXL3 projections in #806, #740 and #741. We plan to keep one GLM family on CUDA, and #804 is the one to build on, since it already has any number of ranks, the prompt cache and the server. The EXL3 path and #839's speed work would come in through your family. We are coordinating in #832 and #593. Looking forward to tonight's tables from |
|
@Bizuayeu thanks for the GLM-5.3-Flash CUDA family. I tested Setup. 4× DGX Spark (GB10), switchless ring with neighbour links only; ranks 0↔2 and 1↔3 are not mutually reachable over the selected RoCE transport. Kernel 7.0.0-1019-nvidia with Default With
These checks compare execution paths within TP=4; they do not establish bitwise equivalence with TP=2. I also saw the rank agreement check reject different checkpoint folder names: all four ranks logged REFUSE naming the snapshot-folder mismatch. That checks the folder name; the LFS content check above was separate. Performance with
At four concurrent streams (
The prefill prompts use After the decode and prefill runs with drafts and The build summaries, with many steps cached, reported Not run: A documentation or startup note naming |
|
@ashhart, thank you. The run of @SvangenStudios, thank you for running After the gate, as separate PRs if that suits:
|
|
@Bizuayeu thanks, that plan sounds good. I'd be happy to test the neighbour-relay split on the four-node ring when it's ready, comparing it with |
|
@ashhart, the description now has every row from |
|
@Bizuayeu, Ling-3.1-flash (#850) needs several pieces this PR already has, and we'd rather extend yours than write second copies. Nothing here asks for changes to this PR. We'd do the extensions as follow-up PRs once it lands, if you agree.
Would you take these as follow-ups on top of this PR, or would you rather keep the GLM kernels family-only and have us copy what Ling needs? |
|
@jschmied, yes, please extend these rather than copy them; follow-ups once #804 lands suit us, as long as GLM's path keeps its bits.
Our own follow-ups after the gate (a startup check that every rank pair can reach each other for |
|
@Bizuayeu, one thing we found while checking Median ms per call, synthetic rows, one GB10 (a conversion was running on it, so absolute times are a little high; the ratio held across runs):
32 heads is one rank of GLM-5.3-Flash at TP=2. Lowering |
|
@jschmied, thanks for timing these. GLM already sends every decode window to the three-kernel chain, a single row and drafted windows alike: For Ling, passing |
…rence, not an f32 row of its own
|
@ashhart, pushed |
4 participants
What this changes
The family #593 tracks: GLM-5.3-Flash on CUDA, tensor parallel over two GB10s (and three, below), on the shared NVFP4
layer (
qlinear,experts.Grouped) and hand-written.cufor the rest. It serves throughtensorfold-native serve(rank 0 serves, the other ranks follow its steps), drafts (the checkpoint's MTP layer) and concurrency go through the
lanes core, and kept prompt states through the family's
snap_*hooks.This branch is 1.0.10 with the pieces split out first (#742, #759 and #761 are in 1.0.9, #800 and #805 in 1.0.10);
until the rest land, its first commits are theirs, and the review is the last five:
now), core: a checkpoint's pages out of the page cache once a loader is done with them #766 page cache release, native: a CUDA family's prompt cache, weight share and open errors come from its own declarations #762 a CUDA family's own prompt-cache and share hooks, cuda: the staged NVFP4 prompt entries leave an expert's items alone (skip), as the decode form does #764 skip on the staged prompt
entries, cuda: Nvfp4Experts.packGateUp packs gate and up in place, no scratch #763 packGateUp in place
swiglu's 64-column block where 512 doesn't divide a rank's width (TP=3's last rank)
Engine.initadmitsconfig.jsonthroughmodelopt.admit; every rank reads its room once its weights are in,and rank 0 admits the streams and the prompt cache by the smallest, naming that rank when it refuses
Next to the Metal GLM
One family package:
config.zigis shared (this addshead_unit,shared_inter,norm_topkand the parse checksthe NVFP4 config needs), both backends sit behind the lane core's backend table, kept prompt states behind
api.prompt_cache.Snapshots, and the tokenizer, chat template and tool parsing stay above the family. The CUDA hostfiles are
cuda_*.zigbeside the Metal ones; they share no code today, for reasons in the code:weights_fix.zig); CUDAreads ModelOpt NVFP4 as stored with bf16 attention projections, cut per rank (
cuda_weights.zig,cuda_shard.zig)mtl.Bufferslices with the position on the host; CUDA one device arena with theposition on the device, which its kernels read
mtl.ComputeEncoder; CUDA launches kernels and NCCL exchanges on a streammirror.zig,cuda_follow.zig), over an MCDMA control window on Metal andncclBroadcaston CUDAmsync/mincoreagainst Linuxmadvise/fadvise(core: a checkpoint's pages out of the page cache once a loader is done with them #766)What could move into shared code in follow-ups, by value: one backend-neutral tensor plan both loaders read (checkpoint
name, dtype and shape, the axis cut, each rank's span; the names and cut axes already agree); a state-size table both
state.zigandcuda_state.zigtake (the arithmetic is the same); the step tape apart from its transport; thetoken-handle ring this family's lanes share with Nemotron's CUDA lanes, into the lane core. Splitting a Config across
machines differs too (Metal rewrites the Config to one Mac's share; CUDA keeps it whole and cuts with
Share, whichallows uneven shares); taking
Shareon Metal would be the larger change, so it is a suggestion here.Several requests in one forward
prefill_step: a prompt goes one 2,048-row chunk at a time and the other streams' rounds run between chunks.hidden_rows: up to 8 streams' verify windows (32 rows) share one forward; row-wise work (projections, experts,exchanges, head) runs over all rows, each stream's KDA chain and DSA attention on its own sequence.
glm-shared-checkholds every stream against a twin that runs each round alone: logits, tokens, drafts and stateequal byte for byte at TP=2 on the full model.
qlinear/experts.Groupeddo.Speed
GB10s in one sitting (median of 2): prompt 8,192 tokens 5.86 -> 5.26 s and 32,768 tokens 22.64 -> 20.54 s, with the
same 49 greedy ids after each prompt; greedy decode without drafts 17.10 to 17.15 -> 17.10 to 17.20 tok/s (three chat
prompts of about 2k tokens).
What we tried for speed during the port, kept and dropped
Measured while porting, each against the build before it in the same sitting: with Nsight Compute on one GB10 (the
five-layer forward check with both TP=2 ranks on it, over a 32k-token prompt where noted), with Nsight Systems on
rank 0 of the two GB10s (a decode step), or end to end on the two. Every kept change leaves each output's bits as
they were.
Kept
_bmmsetup (BLOCK_N 32 or 64, 1 to 3 cp.async stages, 4 warps) and issues the next stage's copy before waiting on this oneDropped
Looked at and left as is: upstream's NVFP4 lane matmul (
qmmf, 1.1% of a decode step, Systems) and the sampled draw(a sampled step 0.1% longer than a greedy one). Neither Nsight Compute (fixed clocks, caches flushed) nor a cache-cold
bench ranked the decode matmul setups as the served decode did, so decode changes were judged on the two GB10s.
Receipt
Environment: this branch at
6380c8daon 1.0.10 (6c5e5192); the rows below are from one sitting on two NVIDIA GB10(driver 580.178.04, CUDA 13.3, NCCL 2.30.7), Linux 7.0.0-1019-nvidia with
kho=off, over RoCE with NCCL;nvidia/GLM-5.3-Flash-NVFP4at423acf37; BF16 KV;--context 40960 --parallel 2for the prompt and decode rows,--parallel 4for the concurrent ones.Exactness:
tools/bench_concurrent.py --alone --serial(levels 1 and 4, both prompts, 256 tokens, temperature 0, 3repetitions): all 30 replies equal their requests alone, and the alone runs equal
"draft": false.resumed = fresh: a chat conversation whose later turns resume the state kept at the history end (cached 2,048 of
about 3,100 tokens) gives the same replies as a server without the cache, greedy, with drafts and without, and
sampled (turns at 3,094 to 3,113 tokens).
In the engine at TP=2:
glm-shared-check(108 rows of 5 streams over 12 rounds equal their lone twins byte forbyte);
glm-lanes-concurrentandglm-lanes-cover(cancels, discards mid-decode and after a rejected draft, eosat 0, 10, 33 and 101 tokens, max_new 1 to 5, a pass resumed at 2,048 and 4,096 against the fresh pass, 512 tokens
with drafts) equal their plain greedy reference; rank 1 follows to the end.
Admission over the ranks (TP=2): with rank 1 alone under
TENSORFOLD_MEMORY_RESERVE_GIB=21,serve --context 40960 --parallel 8stops before serving: rank 0 refuses the eight streams, naming rank 1.Float64 references for every projection whose reduction order is the kernels' own (
tf-cuda-test glm-ref, seed593), the worst column's error over its tolerance:
Every check also holds its kernels' outputs bit-equal across row counts and forms (prompt rows against decode
windows, fused against three-kernel KDA). The bf16 ratios near 1 are the output rounding itself (2^-8 |y|).
Decode speed (
tools/bench_openai.py, 256 tokens, temperature 0, 3 repetitions, median tok/s), with vLLM0.28.1rc1.dev580 (nightly
385dce36) on the same checkpoint and pair, with the profile we serve it with (FP8 KV,MTP k=3, no CUDA graphs, one sequence), in the same sitting:
--no-draftsFour concurrent streams (
bench_concurrent.py, level 4, thechatandcodeprompts): 52.8 and 68.0 tok/s in all,against 27.9 and 33.4 alone.
Prompt speed (
tools/prefill_cold.pyas of394d566c^, which the template names and 1.0.0 removed; itsrunisstandard library only; lengths 2k to 32k, as the server holds 40,960, three prompts each with a unique first line,
a rest after each long one outside the timed span), median TTFT in seconds:
A decode beside a 32k prompt waits at most 1.68 s between tokens (the prompt in 2,048-row chunks).
Memory: the server's
/memorypeak 97.3 GiB of device memory on rank 0 (--context 40960 --parallel 2, throughthe 32k prompts); each host's memory in use above its idle baseline at most 100.5 GiB (this PR) and 111.7 GiB
(vLLM,
gpu_memory_utilization0.80), sampled every second.Tests run (empty cache):
zig build test-cuda-host511/511,zig build test617/622 (5 skipped, none failed),tf-cuda-test quant-ref4 andglm-ref11 passed, the lanes checks on a synthetic two-layer checkpoint thattf-cuda-test glm-synthwrites (cover 40, concurrent 38, shared 1) and on the model cut to five layers (cover 36,concurrent 38); at TP=3 on three GB10s (uneven shares: 22/21/21 heads, 704/704/640 expert rows)
glm-lanes-concurrentand
glm-shared-check;python3 tools/zig/lean_check.pyfinds nothing new against 1.0.10.Not run, and why:
config.zigis shared; its parse and width tests run inzig build test.NCCL_COMM_ID, and NCCL then looks for a bootstrap interface on rank 0'ssubnet, which a ring of /32 addresses doesn't give the other ranks. fabric: tensor parallelism across CUDA processes (bootstrap, NCCL, mailbox and RoCE collectives, plan link, fail-fast) (DeepSeek-V4.1 stack 3/14) #636's startup covers this; the family will
take it rather than carry its own. The engine's TP=3 checks above pass.
EXL3's unit, and cuda: EXL3 projections write fp32 too (qlinear decodeF32 and promptF32) #806 gives EXL3's projections the fp32 outputs
qlinearhas for NVFP4 (cuda: the NVFP4 lane matmul and prompt GEMM write fp32 too (qlinear decodeF32 and promptF32) #757).6380c8daat TP=4 on four GB10sin a switchless ring (a comment on this PR):
glm-lanes-concurrent,glm-shared-checkand the server'sbench_concurrent.py --alone --serialpass withTF_GLM_PREFILL_REDUCE=gatheron every rank. The default splitexchange sends rows straight between every pair of ranks, so a fabric where some pairs cannot connect needs
gather. @BHCC2025's EXL3 glue on this branch passesglm-shared-checkandglm-lanes-concurrenton two andfour GB10s (GLM-5.3-Flash and full GLM-5.3 on the native CUDA engine (tracker) #593).
tools/bench_cache.py --require-reuse: its repeats are one completion prompt with no chat history, andcore/prompt_cache.zigkeeps states at a chat's history end and at shared prefixes, so a CUDA host reports themcold (kept 0); resumed = fresh is shown with a chat conversation above.
Checklist
"draft": false, a resumed prompt equals a fresh one, and concurrent equals solo.glm-refbounds against float64.
vLLM above; the shared kernels it carries leave prompt processing as it was, as their PRs say.)
tests have no code to run before it. Its host tests run in
zig build test-cuda-host, which does no GPU work,and read checked-in fixtures, not the checkpoint, so they have nothing to skip; the GPU checks are
tf-cuda-testcommands.)