Repository navigation
Conversation
… as copied rows The plan split the prompt cache into snapshot slots and copied rows, and a family on pages had to declare a slot's size. A family whose kept states are copies alone (GLM-5.3-Flash on CUDA) has no slot to name: snap_bytes 0 now means no slots, and the whole cache goes to copied rows. On the native host, such a family's cache came from what every stream at a whole window leaves of the room. On pages that count has no meaning and leaves nothing, so the cache is the one the plan already set aside.
The latent write, the pooled keys, the index scores and the dense and sparse latent chunks read a cache's rows through a table of 512-row pages, pages a stride apart, as Nemotron's attention takes pt and the elements between pages. A null table keeps the rows in order, so every caller today passes one and reads the same addresses. The chunk kernels' key loads are made twice from one body and chosen once a call by whether the stream has a table: rows in order in main's loop, or each key's row through the table. A form that branched on the table inside the cp.async loop made nvcc unroll it less even without a table (14 async copies where main has 22), and long prompts on windows ran about 3% slower than main. A scores warp stages 16 pools from a multiple of 16 and a page holds 128, so it reads one page; it, the pooled keys and the latent write look their page up once a block or a stage. Only addresses change: the arithmetic and its order do not. The FP8 latent cache's kernels take no table, and its launchers refuse one (Fp8OnPages).
…'s pages A paged sequence holds its fixed part (KDA states, the window's scratch, the index rings, its first-draft row and a table of 512-row pages) and finds its cache rows through the table. A page holds every DSA layer's and the MTP layer's 512 latent rows and their 128 pooled keys: 13,056 bytes a row on GLM-5.3-Flash. Pages makes the page arena and the paged sequences at once. On rank 0 a stream's pages come from the core pool and go back when it ends; another rank maps the ids rank 0 names. The engine's own sequence keeps its whole window. Paged sequences hold BF16 rows. A kept prompt state takes a paged sequence's rows page by page into the snapshot pages, in the same bytes as a window's prefix, and its page table stays its own, so a state kept on one kind of sequence resumes on the other. A paged sequence holds its own request's rows only, so a prefix's bytes are sized on the engine's own sequence when a paged one is bound. The server makes no pages yet.
…core pool Every stream made a sequence sized for the window on its first bind, and admission priced every stream at that size: at a 1M window, 12.9 GiB a rank whatever the conversation's length. The family now declares a kv_pool.Layout (ashhart#781), as Nemotron does (ashhart#935): a stream's fixed part and 13,056 bytes a token in pages of 512 rows. The server plans the pool beside the engine's own full-window sequence, counted once, and attach makes the paged sequences and the pages. A request takes the pages of its prompt, its reply and one verify window when it starts and gives them back when it ends; admission waits while the pool lacks them. The prompt cache keeps its copied rows beside the pages. TF_KV_PAGES chooses which streams page. auto, the default, pages from three ranks up and keeps a window each below: on two GB10s the room left after the weights (6 to 9 GiB) is under the default 16 GiB prompt cache, which the plan funds after a page a stream and before any other page, so the pool would hold little more than a page a stream. on pages every stream past the engine's own at any rank count, off (or 0, as for Nemotron) gives every stream a window as before, and all puts every stream on pages for checks, the engine's own sequence left idle. TF_KV_POISON=1 fills a stream's pages with NaN before its rows are written, as for Nemotron. On more than one rank, rank 0 plans by the smallest room and sends the pool's size in a new pool step; every rank makes the same pool, and each stream's begin step carries its page ids, which the other ranks write to their tables without taking pages of their own. A page size that differs between the ranks is refused by name. With --kv-cache fp8 the streams keep their windows: auto does not page, and on and all are refused by name (Fp8OnPages), as the FP8 cache's scales have no place on a page yet.
…r byte glm-paged-check runs each prompt on a paged stream and on a twin with a whole window: lengths at a page's edges (511 to 513 rows), at a prompt chunk's (2,047 to 2,049) and past the dense pass (5,000, whose sparse pass reads pooled keys across pages). It compares the first token, its logits and the head's drafts, then windows of 1 to 4 rows across the page edge, their keeps and drafts. A second stream takes and gives back pages between prompts, so the paged stream's pages come out of order, and every page starts as NaN. At the end every page is free again. At TP=2 and TP=3 the other ranks build the pool from rank 0's step and map its ids.
…a host test The pool step's check that a KV page has the same bytes on every rank sat in the follow loop, where no host test reaches it. It moves into Pages.follow, which the loop calls: a follower whose page has other bytes than rank 0's returns RankMismatch before it allocates anything, and with the same bytes it makes the pool as before. The new test drives both on the fake driver. The error log is skipped under test, as the idle report is, since the test runner counts an error log as a failure.
…dense chunk's page once a block, a sparse key's row from the token list On pages the chunk kernels looked each key's page up in a loop of their own, and at TP=2 a paged stream's long prompts ran about 2% slower than main. TF_GLM_PROF put all of it in the MLA stage, 16% slower, with the indexer and the cache writes level. The loop is main's again for every stream: a dense chunk's 512 keys are one page, so a block points its key rows at that page once, and tf_glm_dsa_tokens writes a paged stream's selected tokens as their rows in the pool's pages (a table read a token, once a selection), which the sparse pass reads as main reads the tokens.
… table out of order The six kernels this PR changes run gates 1-4 on a table that is out of order and skips pages, with the cache's part past its page's start as a second cache's is in the pool: tf_glm_lat_write and tf_glm_pool_keys write rows and pooled keys onto pages, tf_glm_scores reads the pooled keys from them, tf_glm_dsa_tokens writes the selected tokens as their rows in the pages, and tf_glm_chunks_16 and tf_glm_chunks_32, dense and sparse, run through tf_glm_merge against a float64 softmax. Scores' sums are exact on integer inputs with power-of-two scales; relu's NaN is a known finding (fmaxf, as Triton's tl.maximum). The pooled keys' and the chunks' sums go through softmax weights, so no input makes them exact in fp32 and they are checked against bounds. docs/kernel-contract.md lists the kernels under Coverage and scores' NaN under Known findings.
….0 and 13.3 on every SM kernel-contract-limits read tf_glm_lat_write, tf_glm_pool_keys, tf_glm_scores, tf_glm_chunks_16, tf_glm_chunks_32 and tf_glm_dsa_tokens from fatbins built for sm_80, 86, 89, 120 and 121. nvcc 13.0's usage goes in glm_cuda_records.cuda, and the four launches nvcc 13.3 compiles differently (lat_write, scores, chunks_32 and dsa_tokens) in a cuda_133 set on 13.3's; nvcc 13.2 was not at hand. Every launch fits every SM. The contract's image list takes glm_ops for tf_glm_dsa_tokens, and the docs name the four in 13.3's set and say the scores' relu is the Triton glue's maximum.
…ize while a paged stream is current Two guards of the paged snapshot code had no test, as breaking each showed: Plan.piece adds the offset inside a page's span when a 2 MiB snapshot page ends inside it, and bytes measures a prefix on the engine's own sequence when the current one is paged (the prompt cache asks for a size before a stream is bound, so the current sequence is whichever ran last). The fixture's state fits one snapshot page, so no copy there starts inside a span: the first test reads piece on three pages out of order. The second makes a sequence of two pages current and asks the size of a 4,096-row prefix, past its pages: it equals the engine's own sequence's.
… slot past the paged sequences waits With TF_KV_PAGES=all the engine's own sequence stays idle, so a server of N lanes has N - 1 paged sequences. Requests that arrive together are admitted before any of them opens (the pool's check sees every paged sequence free), and the last to open took slot N: binding it made a whole-window sequence before the pages refused it with LaneMemory. At TP=3 with a 1M window that was 12.88 GiB on rank 0, which the host's memory guard stopped; at TP=2 a 128k window, 1.77 GiB left behind. The slot's pages now come first, so it waits for a paged sequence and nothing is made for it.
…s git diff --check asks
…em planned before the other streams
…m first and says so at startup
…its pages are a stream's whole table
…s no whole window to measure
…pages serve the other streams while it is short
…, on pages from its init
…, on one GPU and at TP=2 and TP=3
Bizuayeu
marked this pull request as ready for review
October 23, 2026 04:41
This was referenced Oct 23, 2026
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
What this changes
Builds #1163's question 2. With
TF_KV_PAGESpaging (autofrom three ranks up,on), the engine's own sequence (slot 0) takes pages from the core pool too, so no stream holds a whole window. The plan sets one window of pages aside first, so one conversation can still reach--context; while the engine's own conversation is short, those pages serve the other streams. At TP=3 with a 1M window, the pool goes from 1,284 pages beside the engine's 12.88 GiB window to 3,355 pages for all eight streams (receipt below).Stacked on #1163: its 14 commits come first, unchanged, and this pull request's own are the last seven (15 to 21). Until #1163 lands, the diff shows both. If you would rather have these in #1163, I will move them to its branch.
Commits:
62e2aeb3corekv_pool:Budget.engine_paged, false for Nemotron, whose tests pass unchanged. When it is true,plantakes one window of pages (pages(window), the floor) from the room first, then the other streams' fixed parts and a page each, the prompt cache, and the rest as pages, at most(1 + others) × pages(window). A room under the floor isNoStream.holdscounts the same way.7d4f8dc8native:Loaded.engine_pagedgoes toplanandholds, and the startup line and theNoStreammessage say so (below).8737d30dGLM declareswindow_rows(8) to the native server. Without it the plan's window was the context alone, sopages(window)could be one page short of a request's longest page table (16 against 17 at a context of 8,192), and the floor depends on it.b216a74bA kept prefix's size (Plan.size) counts its rows to the capacity, whatever pages are mapped.cuda_snapshots.bytesmeasured a prefix on the engine's own window while a paged stream was current; with no window anywhere it would return the maximum, and the prompt cache would keep nothing, silently.27feda8bThe engine's own sequence on pages:Engine.initmakes slot 0 a paged sequence with no rows of its own;Pagescovers slots 0 to N − 1 and places slot 0 on the arena (State.place), refusing a slot 0 that holds a window (OwnNotPaged); every rank resolvesTF_KV_PAGES, and the ranks agree on it (Settings.pages,Fingerprint.pages, layout 12);allbecomeson, and the lanes loseall's idle slot.8e705835glm-paged-check's paged stream is slot 0, on pages from its init; its whole-window twin stays in slot 3.87e4417eThe raw receipts indocs/receipts/glm-kv-pages-8e705835/. Docs only.What changes for a launch
TF_KV_PAGES, and the ranks agree on it at startup: ranks that resolve it differently stop withTF_KV_PAGESnamed. In glm/cuda: KV pages from the core pool, kept states sized by their position (#1142) #1163 rank 0 alone read it, so a launch that set it on rank 0 only now sets it on every rank.allison: with the engine's own sequence on pages there is no idle slot, and--parallel Nserves N streams on pages.off,autoat one and two ranks, and--kv-cache fp8.#1163's question 1
With the engine's own sequence on pages, the plan funds one window of pages before the prompt cache, the floor question 1 asked about, so
onalways keeps room for one conversation to reach--context. The rest of question 1, whetherautoshould page at one and two ranks, stays open: this pull request keeps windows there.Receipt
On this head (
87e4417e, on #1163's headd7ff6494, on mainf8d254c5). Every run on one GPU and at TP=2 used the binaries built from commit 20 (8e705835); commit 21 adds the raw receipts. TP=3's ranks ran commit 20's engine source with the two commits our ring of three GB10s needs that #1163's receipt describes (tp3/ring.patch). The kernels are #1163's, unchanged, so its kernel contract and gate 5's records stand; the environment is #1163's (GB10s, sm_121, driver 580.178.04, NCCL 2.30.7, Zig 0.17.0, nvcc 13.3.73,nvidia/GLM-5.3-Flash-NVFP4@423acf37).Exactness:
tf-cuda-test glm-paged-checkon the checkpoint's first five layers, its paged stream now the engine's own sequence on pages from its init: at TP=2, at TP=3 and on one GPU, 7 prompts at the pages' edges and 91 window rows equal their whole-window twins' logits, tokens and drafts byte for byte, the pages out of order (12 jumps) and NaN before use, the pool idle after; the other ranks made rank 0's pool and mapped its pages through every step. The synthetic checkpoint's checks pass (glm-ref,glm-synth,glm-lanes-cover,glm-lanes-concurrent,glm-shared-check,glm-paged-check).--context 131072 --parallel 4 --prompt-cache-gib 4 TF_KV_POISON=1:TF_KV_PAGES=allandon, now one configuration ("4 streams at once, the engine's own and 3 more taking 803 pages of 512 rows (5.00 GiB) as they need them (enough for one at a 131072-token window), 180 MiB fixed each"; glm/cuda: KV pages from the core pool, kept states sized by their position (#1142) #1163 had 504 pages with the engine's own sequence idle, or 509 beside its window): 12 cells each, every reply equal to its alone run (28/28) and to"draft": false(30/30), each record's 44 requests cached as planned. Theonrecord replayed on this head's windows (autokeeps them at TP=2) and on mainf8d254c5: against theonand theallrecords, 43/43 replies equal and as cached as planned in all four, drafted == plain 23/23, no resumed reply slower than its fresh one. Images on pages (561 pages): the same tokens drafted, repeated, with"draft": falseand three at once, and a chat resumed past its image (6,144 of 7,603 tokens cached) equals its cold run.--parallel 8,TF_KV_POISON=1on pages:TF_KV_PAGESunset at a 1M window: "8 streams at once, the engine's own and 7 more taking 3355 pages of 512 rows (20.89 GiB) as they need them (enough for one at a 1048576-token window), 131 MiB fixed each, of 37.8 GiB left after a 12.2 GiB reserve" (glm/cuda: KV pages from the core pool, kept states sized by their position (#1142) #1163's: 1,284 pages beside the engine's 12.88 GiB window). 16 cells, every reply equal to its alone run (60/60) and to"draft": false(72/72); the record's 44 requests cached as planned; cold prompts of 1k to 64k in 0.73 s to 43.4 s, as glm/cuda: KV pages from the core pool, kept states sized by their position (#1142) #1163's; seven 2k prompts sent while a 64k conversation held its rows on pages (63,488 of its 65,536 tokens resumed): all seven equal their alone runs, and none waited.TF_KV_PAGES=allat a 96k window: 1,544 pages, the bound of eight windows' pages. 60/60 and 72/72, and in the cells of eight no request waited (the longest first token 1.69 to 2.15 s; under glm/cuda: KV pages from the core pool, kept states sized by their position (#1142) #1163'sallthe eighth waited, 12.0 to 21.2 s).TF_KV_PAGES=off --prompt-cache-gib 0 --parallel autoat a 1M window: two windows, as before. Theautorecord replayed there: 43/43, drafted == plain 23/23.TF_KV_PAGES=allat a 96k window, asking a 37.2 GiB prompt cache. The plan trims the cache to 35.9 GiB to keep the floor and a page a stream, 200 pages. Three 64k prompts (129 pages each) alone, then together: all three equal their alone runs, each resumed 63,488 of its 65,536 tokens, and they ran one at a time.free N + store 0 == total N: 3,355, 1,544 and 200 at TP=3; 803, 757 and 561 at TP=2.Speed at TP=2 against main
f8d254c5on the same two GB10s, the serves alternating: this head, main, this head, main, then this head with every stream on pages (TF_KV_PAGES=all), each at--context 131072 --parallel 4 --prompt-cache-gib 0. Seconds, lower is better; a ratio is this head's faster serve over main's slower.fibonacci-raw, s/tokengpu-chat-no-think, s/tokenAt two ranks
autokeeps windows, so "this head" serves windows; "all on pages" runs every stream on pages, a lone request on the engine's own paged sequence. Greedy decode and the 32k chat gave the same tokens in all five serves. At 64k main's own two serves differ by 0.7%.Checks on a GB10 (nvcc 13.3.73,
-Dsm=121) on commit 20: the build,test-cuda-host(2,516 of 2,522 tests, 6 skipped),test-kv-pool17/17 andtest(276/276 steps) pass, andzig fmt --checkon the changed files is clean.The no-GPU CI's steps under WSL2 on commit 20:
zig build test3,470/3,486 (15 skipped; onlycuda_cost_identityfails, to WSL2's coarse file times as before),test-kernel-contract72/72,test-golden, the server and pull security tests 20/20,fuzz-cortheon14/14, the served Flash Next unittest (8), and the GLM snapshot, KV pool and CUDA host tests (72/73, the same failure). lean: 438 problems, as on main, 0 new.Thirty mutations under WSL2, each guard broken alone on a copy of commit 20, with
test-kv-pool,test-cuda-hostandtest-glm-snapshots-host: glm/cuda: KV pages from the core pool, kept states sized by their position (#1142) #1163's fifteen and fifteen new ones. 29 fail a test, and the unchanged copy fails none; glm/cuda: KV pages from the core pool, kept states sized by their position (#1142) #1163'sbytes-on-paged-seqno longer applies, as its line becamePlan.size, which two of the new ones break. The new fifteen:ownPages0)kv_pool_test"the engine's own sequence on pages: a window of them comes first, ..." andcuda_kv_pool"kv pages: an engine whose own rows are pages plans a window of them first, ..."holdswithout the floorplannot told the engine's rows are pagescuda_kv_pool"kv pages: an engine whose own rows are pages ..."holdsnot toldwindow_rows0cuda_native"the plan's window, the context and a verify window, takes the pages of a stream's table at that capacity"glm_snapshot_host"while a paged stream is current, a prefix's size is the engine's own sequence's, past its pages" and "with the engine's own sequence on pages too, a prefix's size past a stream's pages is a window's"cuda_pages_test"TF_KV_PAGES: a stream past the paged sequences waits for one, ..." and two morecuda_pages_test"the engine's own sequence holds no rows before its pool, then sits on the pool's pages ..."cuda_pages_test"a pool is refused while the engine's own sequence holds a window or was never made, ..."cuda_pages_test"TF_KV_PAGES: a stream past the paged sequences waits for one, ..." and two morecuda_pages_test"TF_KV_PAGES: the lane host admits a request by the pool alone, the engine's own sequence free or not"pagesleft out of the ranks' agreed settingscuda_ranks"every scheduling and identity field that differs stops the ranks" and "fingerprints: ..."pagesleft out of the fingerprintcuda_ranks"fingerprints: agreeing ranks pass; the first differing field and rank are named"Not run, and why:
--drafter): no DFlash2 weights on these hosts.Checklist
"draft": false, a resumed prompt equals a fresh one, and concurrent equals solo. (At TP=2 and TP=3, above.)SUMMARYline is in the receipt. (No kernel changes; glm/cuda: KV pages from the core pool, kept states sized by their position (#1142) #1163's registrations andSUMMARYstand.)