Repository navigation
glm cuda: a prompt chunk's indexer rows split across the ranks, their picks all-gathered - #1404
Conversation
|
Updated the description: it now opens with the measured effect on my hardware, prompt and decode before and after against main, taken from the receipt below. No code change. |
|
Stacked with #1423, a paired step silently drops this PR's indexer split, and |
|
Thanks, Bizuayeu. This is a clean change, and the exactness argument holds up on reading:
On main For the record:
We plan to land this before #1423. #1423 will then take the fix you described in your comment: one split choice for |
…and counts it once (ashhart#1404) Engine.prefillChunk reads and clears the split flag before it hands a paired step to prompt_pair.prefill, which never set it on its two forwards, so with ashhart#1404 a paired step ran the indexer whole on every rank. The step's flag now reaches both chunks: the header carries the one choice rank 0 makes at the step's start, and index_splits counts the step once. The twin buffers already share the first chunk's select scratch, whose pools hold ashhart#1404's world rows more.
… tf_glm_scores and tf_glm_select_rows (TF_GLM_INDEX_CUDA) The kernels are ported from BadAd84's Python patches 0101 (one-pass top-K pools) and 0105 (pool scores), for #1430. tf_glm_index_scores keeps tf_glm_scores' arithmetic exactly (the mma chain over the 8 k16 steps from zero, the relu, the head sum and its shuffle order) and changes only the data movement: a CTA of 8 rows shares each staged 64-pool key tile instead of every row staging its own. tf_glm_index_pools replaces tf_glm_select_rows for a prompt block: a strided-sample floor, one order-preserving compaction into shared memory, then the same radix select and ascending write over the candidates, or over the whole row when the candidates miss [512, cap]; the candidate set holds the 512 best and every tie at the 512th, so the pools are the select's, ties to the lower pool. TF_GLM_INDEX_CUDA=1 takes both for chunks of 64 or more rows, =scores only the scores; unset, nothing changes. A build without the glm_index fatbin logs and keeps the existing launches. Decode and #1404/#1424's areas are untouched; the hook is two branches in selectTokensProfile and an optional field of glue.Set. Host: tools/glm_index_emulate runs both kernels as host C++ (real barriers, warp exchanges, ldmatrix and mma models) against plain models, in test-glm-index-host and test-cuda-host. GPU: tf-cuda-test glm-index-check compares selectTokens with and without the kernels byte for byte at 1k-262k pools, plus tied scores.
|
Our GB10 pair run on the stack (this PR with #1352 on top) is done: GLM-5.3-Flash NVFP4 at TP=2, main against the stack, alternating main, stack, stack, main, at a 163,840-token window.
Since then #1449's CUDA indexer has landed in the same function. The stack rebases onto main without conflicts, and we've done that on our side. Before landing we'll run one more pair cell with |
…ave room under the line limit
…exer rows split across the ranks
…, the first position with sparse rows
d287834 to
0bc0d3d
Compare
|
Thanks @Bizuayeu, this is merged, together with #1352. From our runs on two GB10s at TP=2, alternating main and this branch:
Our full test (the Linux cross-compile, release paths and |
…and counts it once (ashhart#1404) Engine.prefillChunk reads and clears the split flag before it hands a paired step to prompt_pair.prefill, which never set it on its two forwards, so with ashhart#1404 a paired step ran the indexer whole on every rank. The step's flag now reaches both chunks: the header carries the one choice rank 0 makes at the step's start, and index_splits counts the step once. The twin buffers already share the first chunk's select scratch, whose pools hold ashhart#1404's world rows more.
|
Thanks @Bizuayeu: the indexer split across ranks shipped in TensorFold 1.1.9. |
…and counts it once (ashhart#1404) Engine.prefillChunk reads and clears the split flag before it hands a paired step to prompt_pair.prefill, which never set it on its two forwards, so with ashhart#1404 a paired step ran the indexer whole on every rank. The step's flag now reaches both chunks: the header carries the one choice rank 0 makes at the step's start, and index_splits counts the step once. The twin buffers already share the first chunk's select scratch, whose pools hold ashhart#1404's world rows more.
…and counts it once (ashhart#1404) Engine.prefillChunk reads and clears the split flag before it hands a paired step to prompt_pair.prefill, which never set it on its two forwards, so with ashhart#1404 a paired step ran the indexer whole on every rank. The step's flag now reaches both chunks: the header carries the one choice rank 0 makes at the step's start, and index_splits counts the step once. The twin buffers already share the first chunk's select scratch, whose pools hold ashhart#1404's world rows more.
…and counts it once (ashhart#1404) Engine.prefillChunk reads and clears the split flag before it hands a paired step to prompt_pair.prefill, which never set it on its two forwards, so with ashhart#1404 a paired step ran the indexer whole on every rank. The step's flag now reaches both chunks: the header carries the one choice rank 0 makes at the step's start, and index_splits counts the step once. The twin buffers already share the first chunk's select scratch, whose pools hold ashhart#1404's world rows more.
What this changes
Each rank now runs the DSA indexer for its share of a prompt chunk's rows, so long prompts get faster; decode and the
output bytes do not change. On my hardware (GLM-5.3-Flash XSFP4,
--parallel 4, the mean coldttft_s):1b4f2f979→ this headfibonacci-rawDecode never splits: at TP=2 this binary's launches with the split off decoded the same, so that 0.3 to 0.4% is the
build's. Details in the receipt.
Every rank runs GLM-5.3-Flash's DSA indexer for every row of a prompt chunk:
tf_glm_scoresover each row's visiblepools, then
tf_glm_select_rowsfor its 512 pools, beforetf_glm_dsa_tokensturns the picks into token lists. Thescores grow with the position (each row reads all its pools), so on three GB10 at TP=3 the indexer was 1.6% of a 32k
prompt and 7.4% of a 128k one (the engine's prompt profile,
index_scoreandindex_select), and it keeps growingtoward the 1M window.
With
TF_GLM_INDEX_SPLITset to a position, rank 0 has each prompt chunk from that position on split its indexerrows: rank r picks for rows
[r·per, min(rows, (r+1)·per))withper = cdiv(rows, world), block by block as one passwould (its
npandloopfrom the whole chunk), and an in-placencclAllGatherof the picks (8 bytes a pool, 512 arow) gives every rank every row's; the token lists are then assembled for all rows as before. Rank 0 carries its
choice in the chunk's header (flag 2), so the other ranks need no setting of their own.
computes them.
tf_glm_scorescomputes each pool's score from its row and pool alone (the block's grouping of poolcolumns,
loopandnb, only decides which block computes it), andtf_glm_select_rowsworks a row at a time.tf-cuda-test glm-split-checkcompares split prompts with whole twins: the first token, the next window's logitsand tokens, and every byte of the state.
waited for the previous site's side-stream exchanges (
gluePieces), so it is issued after them, in the same hostorder on every rank.
poolsgrow byworldrows (the last rank's equal block may end past the chunk).faster at every size measured, and no slower at 32k (below).
TF_GLM_INDEX_SPLIT=0is main's path. The engine's ownin-process ranks (
.local), the MTP head's DSA layer, decode windows andprefillManynever split. The ranks' testsmove to
cuda_ranks_test.zigfirst (a pure move), so the setting fits the line limit. That first commit is sharedwith a GLM pull request I am still measuring (paired prompt chunks); whichever lands first, the other rebases.
Receipt
Environment: TensorFold
e7daae854(the head before its default's commit) against main1b4f2f979; three GB10 (sm_121, 128 GB unified memory) at TP=3over a switchless RoCE ring, and two at TP=2; driver 580.178.04; nvcc 13.3.1; GLM-5.3-Flash XSFP4
(
Bizuayeu/GLM-5.3-Flash-XSFP4at88736f9). Each served with--parallel 4 --prompt-cache-gib 0, launched in theorder main, split from 2048, whole, split from 32768, whole, split from 2048, main (the same binary for every split
setting).
Exactness:
glm-split-checkpasses at TP=3 and at TP=2 ("prompts of { 6444, 4098, 4101, 4096 } rows equal theirwhole twins in the first token, the next window's logits and tokens, and every state byte"; 4098 leaves a rank with
no row at TP=3).
token_sha256is the same over every launch of both windows, for both decode prompts and all threephases of the 32k and 128k prompts.
Prompt speed before and after: the cold
ttft_sfrombench_cache.py --require-cold --streams 1(first, repeatedand concurrent are all cold at
--prompt-cache-gib 0), the mean of six requests (three for 32,768), in seconds:Gains are against whole on the same binary. The long prompt ran on the TP=2 pair only, whole, split and whole
again (three launches, the six whole requests 464.27 to 467.66 s, the three split 422.74 to 422.94 s), and its
token_sha256was the same in all three. Every split request beat every whole one at 128k on both TPs and at 32kon TP=3; at 32k on TP=2 the ranges overlap. The indexer was 7.4% of a 128k prompt at TP=3 in the engine's profile,
so splitting it three ways could save about 4.9%; the window measured 4.6 to 4.8%.
Decode speed before and after (
bench_openai.py --tokens 256 --reps 3 --temperatures 0, medians,fibonacci-raw):at TP=3 main 57.76 and 57.52 tok/s, this binary 57.43 to 57.66 over its five launches; at TP=2 main 44.55 and 44.59,
this binary 44.37 to 44.41 (0.3 to 0.4% below main in that window, whole launches included, so not the split's
doing; decode never splits).
Tests run:
zig build test4,226 of 4,244 passed, 17 skipped, 1 failed (cuda_cost_identity, which fails on mainthe same way on that host's coarse file times),
zig build compile-config-server7/7,zig build compile-cuda-tests10/10, the GLM host steps (test-glm-prompt-settings-host5/5,test-glm-tp3-host7/7,test-glm-pages-host12/12,test-glm-fp8-kv-host13/13,test-glm-protocol-host8/8),zig fmt --check, thelean check (no new problems against main).
Not run, and why: Metal and HIP, which this path does not reach.
Checklist
computes a prompt row's picks changes.)
"draft": false, a resumed prompt equals a fresh one, and concurrent equals solo. (Thebytes are main's:
glm-split-checkand the hashes above.)changes;
glm-split-checkcompares the bytes.)SUMMARYline is in the receipt. (No kernel changes.)1b4f2f979: faster at 128k on both TPs andat 516k on TP=2, no slower at 32k.)
splitAt, the chunk header's flag 2 andSettings.index_split, which main does not have;glm-split-checkneeds aGPU and a checkpoint.)