Skip to content

glm cuda: a prompt chunk's indexer rows split across the ranks, their picks all-gathered - #1404

Merged
ashhart merged 3 commits into
ashhart:mainfrom
Bizuayeu:pr/glm-index-split
Oct 28, 2026
Merged

ashhart merged 3 commits into
ashhart:mainfrom
Bizuayeu:pr/glm-index-split

Conversation

@Bizuayeu

@Bizuayeu Bizuayeu commented Oct 26, 2026 •

Copy link
Copy Markdown
Contributor

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 cold ttft_s):

TP=3, three GB10: main 1b4f2f979 → this head TP=2, two GB10: main → this head
32k-token prompt 19.794 → 19.653 s (+0.7%) 23.886 → 23.754 s (+0.6%)
128k-token prompt 83.958 → 80.257 s (+4.6%) 99.605 → 97.243 s (+2.4%)
516,096-token prompt 465.889 → 422.848 s (+10.2%, against this binary with the split off)
Decode, fibonacci-raw 57.76, 57.52 → 57.43 to 57.66 tok/s 44.55, 44.59 → 44.37 to 44.41 tok/s

Decode 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_scores over each row's visible
pools, then tf_glm_select_rows for its 512 pools, before tf_glm_dsa_tokens turns the picks into token lists. The
scores 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_score and index_select), and it keeps growing
toward the 1M window.

With TF_GLM_INDEX_SPLIT set to a position, rank 0 has each prompt chunk from that position on split its indexer
rows: rank r picks for rows [r·per, min(rows, (r+1)·per)) with per = cdiv(rows, world), block by block as one pass
would (its np and loop from the whole chunk), and an in-place ncclAllGather of the picks (8 bytes a pool, 512 a
row) 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.

  • Why the bits cannot change: a row's scores and picks come from the same kernels on the same inputs whichever rank
    computes them. tf_glm_scores computes each pool's score from its row and pool alone (the block's grouping of pool
    columns, loop and nb, only decides which block computes it), and tf_glm_select_rows works a row at a time.
    tf-cuda-test glm-split-check compares split prompts with whole twins: the first token, the next window's logits
    and tokens, and every byte of the state.
  • The gather runs on the main stream inside the chunk's DSA layers. On the pieced path the main stream has already
    waited for the previous site's side-stream exchanges (gluePieces), so it is issued after them, in the same host
    order on every rank.
  • What it holds: the prompt buffers' pools grow by world rows (the last rank's equal block may end past the chunk).
  • The default is 2,048, the first position with sparse rows: every chunk past it splits. It was as fast as 32,768 or
    faster at every size measured, and no slower at 32k (below). TF_GLM_INDEX_SPLIT=0 is main's path. The engine's own
    in-process ranks (.local), the MTP head's DSA layer, decode windows and prefillMany never split. The ranks' tests
    move to cuda_ranks_test.zig first (a pure move), so the setting fits the line limit. That first commit is shared
    with 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 main 1b4f2f979; three GB10 (sm_121, 128 GB unified memory) at TP=3
    over 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-XSFP4 at 88736f9). Each served with --parallel 4 --prompt-cache-gib 0, launched in the
    order main, split from 2048, whole, split from 32768, whole, split from 2048, main (the same binary for every split
    setting).

  • Exactness: glm-split-check passes at TP=3 and at TP=2 ("prompts of { 6444, 4098, 4101, 4096 } rows equal their
    whole 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_sha256 is the same over every launch of both windows, for both decode prompts and all three
    phases of the 32k and 128k prompts.

  • Prompt speed before and after: the cold ttft_s from bench_cache.py --require-cold --streams 1 (first, repeated
    and concurrent are all cold at --prompt-cache-gib 0), the mean of six requests (three for 32,768), in seconds:

    TP, prompt main whole (same binary) split from 2048 split from 32768
    3, 32k 19.794 19.778 19.653 (+0.64%) (whole: the last chunk starts at 30,720)
    3, 128k 83.958 84.104 80.257 (+4.79%) 80.405 (+4.60%)
    2, 32k 23.886 23.817 23.754 (+0.27%) (whole)
    2, 128k 99.605 99.828 97.243 (+2.66%) 97.326 (+2.57%)
    2, 516,096 tokens — 465.889 422.848 (+10.18%)

    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_sha256 was the same in all three. Every split request beat every whole one at 128k on both TPs and at 32k
    on 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 test 4,226 of 4,244 passed, 17 skipped, 1 failed (cuda_cost_identity, which fails on main
    the same way on that host's coarse file times), zig build compile-config-server 7/7, zig build compile-cuda-tests 10/10, the GLM host steps (test-glm-prompt-settings-host 5/5, test-glm-tp3-host 7/7,
    test-glm-pages-host 12/12, test-glm-fp8-kv-host 13/13, test-glm-protocol-host 8/8), zig fmt --check, the
    lean check (no new problems against main).

  • Not run, and why: Metal and HIP, which this path does not reach.

Checklist

  • Every token still goes through the lane rounds. No serial path, and nothing that needs drafts off. (Only who
    computes a prompt row's picks changes.)
  • Drafted output equals "draft": false, a resumed prompt equals a fresh one, and concurrent equals solo. (The
    bytes are main's: glm-split-check and the hashes above.)
  • No precision traded for speed. If the bits change, the description says which sums change and why. (No sum
    changes; glm-split-check compares the bytes.)
  • A new or changed GPU kernel is registered with the kernel contract (docs/kernel-contract.md); each backend's
    SUMMARY line is in the receipt. (No kernel changes.)
  • Prompt processing is no slower than the last release. (Against main 1b4f2f979: faster at 128k on both TPs and
    at 516k on TP=2, no slower at 32k.)
  • The description says which platforms I ran, Metal M1 to M5 and CUDA, and which I could not.
  • New tests fail before the change, pass after it, and skip cleanly without their dependency. (The host tests read
    splitAt, the chunk header's flag 2 and Settings.index_split, which main does not have; glm-split-check needs a
    GPU and a checkpoint.)
  • Comments and docstrings are one line. No measurements or history in the source.
  • No personal data, machine names, internal hosts or local paths. No AI attribution lines.

@Bizuayeu

Copy link
Copy Markdown
Contributor Author

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.

@Bizuayeu

Copy link
Copy Markdown
Contributor Author

Stacked with #1423, a paired step silently drops this PR's indexer split, and glm-split-check no longer compiles against #1423's Leader.chunk. The fix for whichever of the two lands second, and the stacked numbers (with #1424 as well: a cold 128k prompt 13.0% faster at TP=3, the same tokens), are in #1423 (comment).

@ashhart

ashhart commented Oct 27, 2026

Copy link
Copy Markdown
Owner

Thanks, Bizuayeu. This is a clean change, and the exactness argument holds up on reading:

  • tf_glm_scores computes each (row, pool) score from that row's position and the pool alone, and its 16-pool tiles
    sit at fixed pool offsets whatever the grid. A block that starts at rank·per therefore gives the same bits as the
    whole chunk's 512-row blocks.
  • tf_glm_select_rows works one row at a time.
  • The in-place all-gather puts every row's picks back where one pass would have written them.

On main b30d5b9d3, fmt and lean are clean (no new findings), zig build test and the GLM host steps pass, and the
Linux CUDA cross-compiles build.

For the record: d28783431 makes the split the default from 2,048, so serving without TF_GLM_INDEX_SPLIT now takes
the split path at TP≥2. Our earlier TP=2 run on NVFP4 was this head with the variable unset (equal tokens on all 52
streams), so that was the split path. Before merging we'll add:

  • cold first-token times at 8k, 32k and 128k with the default, with TF_GLM_INDEX_SPLIT=0, and on main (8k because
    every chunk from the second now pays a gather);
  • glm-split-check at TP=2 on NVFP4.

We plan to land this before #1423. #1423 will then take the fix you described in your comment: one split choice for
both chunks of a paired step, and glm-split-check on the new Leader.chunk. Thanks for working that out ahead of
time.

Bizuayeu added a commit to Bizuayeu/TensorFold that referenced this pull request Oct 27, 2026
…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.
ashhart pushed a commit that referenced this pull request Oct 27, 2026
… 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.
@ashhart

ashhart commented Oct 28, 2026

Copy link
Copy Markdown
Owner

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.

  • Tokens: every cell matches main, 36 of 36.
  • Prefill: 2.2% faster at 128k tokens in both of the stack's runs (1,387 to 1,389 tok/s against main's 1,353 to 1,362). At 8k and 32k it's within the drift between runs.

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 TF_GLM_INDEX_CUDA=1, so the split is checked over #1449's kernels too. Then both PRs land together. Thank you, @Bizuayeu.

@ashhart
ashhart force-pushed the pr/glm-index-split branch from d287834 to 0bc0d3d Compare October 28, 2026 15:03
@ashhart
ashhart merged commit 897d0c0 into ashhart:main Oct 28, 2026
4 of 9 checks passed
ashhart added a commit that referenced this pull request Oct 28, 2026
…'s wide decode select (#1424), its indexer split across ranks (#1404) and its share-billed prompt step ending on the GPU (#1352)
@ashhart

ashhart commented Oct 28, 2026

Copy link
Copy Markdown
Owner

Thanks @Bizuayeu, this is merged, together with #1352. From our runs on two GB10s at TP=2, alternating main and this branch:

  • Default settings: the same tokens as main in all 36 cells, and prefill at 128k tokens about 2% faster.
  • With TF_GLM_INDEX_CUDA=1: again 36 of 36 cells equal. Main's first run was disturbed, so against its clean run prefill is the same at 8k and 32k and 1.1% faster at 128k.

Our full test (the Linux cross-compile, release paths and zig build test) and our local gate on main with both PRs pass. Thanks for keeping the decode path out of the split.

Bizuayeu added a commit to Bizuayeu/TensorFold that referenced this pull request Oct 29, 2026
…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.
@ashhart

ashhart commented Oct 29, 2026

Copy link
Copy Markdown
Owner

Thanks @Bizuayeu: the indexer split across ranks shipped in TensorFold 1.1.9.

Bizuayeu added a commit to Bizuayeu/TensorFold that referenced this pull request Oct 31, 2026
…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.
Bizuayeu added a commit to Bizuayeu/TensorFold that referenced this pull request Oct 31, 2026
…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.
Bizuayeu added a commit to Bizuayeu/TensorFold that referenced this pull request Nov 1, 2026
…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.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants