Repository navigation
Conversation
|
Pushed |
… 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.
1dda6c9 to
ae8a5a6
Compare
|
Rebased on main after 1.1.2 ( |
…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.
|
Pushed |
… 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.
|
Pushed |
….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.
|
Pushed |
… 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.
|
Pushed |
…s git diff --check asks
|
Pushed |
|
Question 2, the engine's own window on pages: #1211 builds it, stacked on this pull request (its own commits are the last seven). With pages chosen, the engine's own sequence takes pages too, one window of them planned first so that a conversation can still reach |
What this changes
Implements #1142. GLM-5.3-Flash on CUDA gave every stream a window of KV, 13,056 bytes a token on every rank, so a 1M window held 12.75 GiB a rank for any conversation, and every kept prompt state copied a whole window. As proposed in #1142, the streams past the engine's own now take 512-row pages from the core
kv_pool(#781), as Nemotron's do since #935. A kept state takes only the rows up to its position (piece 1 of #1142), which #993 brought to every stream in 1.1.2; a paged stream's rows go into #993's snapshot pages page by page.TF_KV_PAGES=auto, the default, pages from three ranks up and keeps a window per stream at one and two ranks. That differs from Nemotron's default; question 1 below says why and asks about it.Rebased on main after 1.1.2 (
f8d254c5). The receipts below ran commit 11, and the raw ones are indocs/receipts/glm-kv-pages-0b094b74/.Commits:
2671de38core:kv_pool.splitwithsnap_bytes == 0means a family without snapshot slots, whose whole prompt cache is copied rows;cuda_kv_pool.kept()then uses the paged plan'scopied. Without it the cache came from what window-sized streams leave, which on pages is nothing.2ee34d6btf_glm_lat_write,tf_glm_pool_keys,tf_glm_scoresandtf_glm_chunks_16/32take a page table and the page stride, as Nemotron's attention takespt. A null table reads the contiguous rows, as the engine's own sequence does. The chunk kernels load a window's keys in main's loop and a paged stream's keys in a loop of their own, chosen once a call: a helper that returned a pointer and branched on the table inside thecp.asyncloop made nvcc unroll that loop less even without a table. The other kernels look a page up once a block or a stage. The FP8 cache's kernels take no table, and their launchers refuse one.ff307818A paged sequence (its fixed part, a page table, no cache rows of its own) andPages: the page arena, the paged sequences, leases, table uploads, the NaN fill and the idle census. GLM-5.3-Flash TP=2 (0.6.0 lineage): with uptime a 13-token prefill reaches 18.8 s and decode falls 40 to 14 tok/s, an idle rank 1 reads 96% GPU, and a restart clears both #993's snapshot plan takes a paged sequence's rows page by page (below). Host tests on a fake driver.804c4da3The family'skv_pool.Layout,attach, admission,TF_KV_PAGESandTF_KV_POISON, and the ranks (apoolstep, page ids on thebeginstep). With--kv-cache fp8the streams keep their windows (below).8e279c85tf-cuda-test glm-paged-check: a paged stream against a whole-window twin, byte for byte.ae8a5a62The other ranks make rank 0's pool throughPages.follow, which refuses a page size other than rank 0's (RankMismatch) before it allocates anything; a host test drives it on the fake driver.b7fecb3dOn pages too the chunk kernels read keys in main's loop: a dense chunk's 512 keys are one page, so a block points its key rows at that page once, andtf_glm_dsa_tokenswrites 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. With commit 2's own loop for pages, long prompts on pages ran 2.0 to 2.5% slower than main, all of it in the MLA stage byTF_GLM_PROF.786fabe8The six kernels under the kernel contract, on a page 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_writeandtf_glm_pool_keyswrite onto pages,tf_glm_scoresreads the pooled keys from them,tf_glm_dsa_tokenswrites the selected tokens' rows, andtf_glm_chunks_16andtf_glm_chunks_32, dense and sparse, run throughtf_glm_mergeagainst a float64 softmax: eight registrations.docs/kernel-contract.mdlists them, andtf_glm_scores' NaN under Known findings (below). Tests and docs only.c38f6d0bGate 5's records for the six launches: nvcc 13.0's usage, read from every SM's image (sm_80, 86, 89, 120 and 121), inglm_cuda_records.cuda, and nvcc 13.3's where it differs (four launches) in acuda_133set the registry adds to 13.3's; the contract's image list takesglm_opsfortf_glm_dsa_tokens. Records and docs only.3b41fd7cTests for two guards of the paged snapshot code that had none, found by breaking each (below): a paged part's piece that starts inside a page's span keeps its offset there, and while a paged stream is current a prefix's size is the engine's own sequence's. Tests only.0b094b74A stream takes its pages before its sequence is bound. WithTF_KV_PAGES=allthe engine's own sequence stays idle, so--parallel Nhas N - 1 paged sequences. Requests that arrive together are admitted before any of them opens, and the last to open took slot N: binding it made a whole-window sequence on rank 0 (12.88 GiB at a 1M window) before the pages refused it withLaneMemory. Now that slot waits for a paged sequence and nothing is made for it. Found by the first TP=3 run (below); at TP=2 the same path made a 1.77 GiB window inside the reserve, so the earlier TP=2 replies were equal. With a host test.cbe29802That commit's comment on one line, as the lean check asks.2ce601ceThe raw receipts indocs/receipts/glm-kv-pages-0b094b74/: the logs and JSON of the runs below, with machine names, internal addresses and local paths replaced, and TP=3'stp3/ring.patch(below). Docs only.d7ff6494Two stack traces in that receipt's mutation logs indented with spaces, as the desktop job'sgit diff --checkasks. Docs only.On main after 1.1.2
snap_bytessizes a prefix on the engine's own sequence while a paged one is bound. A host test intest-glm-snapshots-hostkeeps a window's prefix, resumes it on a paged sequence whose pages are out of order, and back, and checks every field and row against the two layouts; two more check a piece that starts inside a page's span, which that test's small state never makes, and the size asked while a paged sequence is current.--kv-cache fp8keeps a window per stream: with itautodoes not page, andonandallare refused by name (Fp8OnPages). Question 3 below asks how you would like FP8 rows on pages.poolfollows itsidleand GLM-5.3-Flash TP=2 (0.6.0 lineage): with uptime a 13-token prefill reaches 18.8 s and decode falls 40 to 14 tok/s, an idle rank 1 reads 96% GPU, and a restart clears both #993'ssnapshotsincuda_follow.zig'sKind. glm/cuda: load and serve nvidia/GLM-5.3-Flash-FP4-Local as stored #1050 is open and touchescuda_dsa_block.zig,cuda_engine.zig,cuda_native.zigandcuda_state.zigin other hunks; whichever lands second is rebased on the first. Both also load theglm_gluemodule incontract_tests.zig'srun, in hunks git can merge without a conflict, so the second to land uses the first's module rather than a secondvar glm_glue; and both addglm_gluelaunch records toglm_cuda_records.zig.Against #1142
capacity/4 + 2rows: the two extra rows arepool_keys' launch slack (gridrows/4 + 2). The kernel writes complete pools only (4p + 3 <= pos + rows - 1) andscoresreads visible pools only (p < (pos + r + 1)/4), both insidecapacity/4, so pages carry no slack.selectTokens'np_max_capstayscapacity/4on pages.Layout
rof a cache is atpt[r / 512] * page + (r % 512) * 1,024from the cache's base in the page, pooled keypatpt[p / 128] * page + (p % 128) * 256. A dense chunk is 512 keys, one page, so a block looks its page up once. Ascoreswarp stages 16 pools from a multiple of 16 and a page holds 128, so it stays in one page. A sparse key's row comes from the token list, wheretf_glm_dsa_tokenswrites it as a row in the pages (commit 7).Admission and ranks
attach: rank 0 plans the pool by the smallest rank's room and sends apoolstep (streams, pages, page bytes). Every rank makes the same paged sequences and page arena; a page size that differs between ranks is refused by name (RankMismatch).ceil((prompt + max_tokens + 16) / 512)pages when it starts and gives them back when it ends. Its page ids ride itsbeginstep; the other ranks write them to their tables and never take pages of their own.fits: the engine's own sequence is free, or a paged sequence is free and the pool holds the request's pages. Otherwise the lane host puts the request back (LaneMemory) until a stream ends.LaneMemory) with no sequence made for its slot.kv pool idle: promised 0, free N + store 0 == total N.TF_KV_PAGESauto(or unset):onfrom three ranks up,offat one and two.on(any value but the others, as Nemotron readsTF_KV_PAGES): every stream past the engine's own on pages.offor0: a window per stream, as before (0is Nemotron's word for it).all: every stream on pages, the engine's own sequence left idle, for checks (one stream fewer at once than--parallel).TF_KV_POISON=1: a stream's pages are NaN until its rows are written.Rank 0 alone reads
TF_KV_PAGES; the other ranks follow itspoolstep. Each rank readsTF_KV_POISONfor its own pages. With--kv-cache fp8,autoandoffgive windows, andonandallare refused (Fp8OnPages).Questions
kv_pool.planfunds a page a stream, then the prompt cache, then the other pages. On two GB10s the room left after the weights and the engine's own window is 6 to 9 GiB, under the default 16 GiB cache, so withonthe pool would be three pages. A request reserves the pages of its prompt and itsmax_tokens, so most concurrent requests would wait for the engine's own sequence: fewer streams at once than with windows. That is whyautokeeps windows at TP=2, and why the TP=2 receipts below set--prompt-cache-gib 4. Would you rather the plan fund pages before the cache up to a floor (the pages of one request at--context, or a share of the room)? Thenoncould be the default at every rank count, as for Nemotron.--context, the peak stays the same; what changes is that those bytes serve the other streams while the engine's own conversation is short. It needskv_pool.planto count the engine's own sequence on pages (aBudgetfield; Nemotron's behavior unchanged) and slot 0 to become a paged sequence. If you agree, we add it to this pull request; if not, it stays as KV pages: one core owner with reservation leases, Nemotron CUDA pilot (#781) #935 has it. glm/cuda: the engine's own sequence on KV pages too, one window of them planned first (stacked on #1163) #1211 builds this, stacked on this pull request, so you can see it before you decide.--kv-cache fp8a cache keeps 512 E4M3 bytes a row and, after all its rows, one FP32 scale a row. On pages the rows would go 512 to a page as BF16 rows do, and a row's scale is easiest found beside its own page's rows: each 512-row block would keep its scales after its rows. Would you have windows keep the same blocks, so that a window is pages in order and one address formula serves both? Until you say, this pull request keeps FP8 caches on windows; FP8 on pages would follow with its kernels' registration.Receipt
On this head (
d7ff6494, on mainf8d254c5). Every run on one GPU and at TP=2 used the binaries built from commit 11 (0b094b74); commits 12 to 14 put a comment on one line and add the raw receipts. The kernel contract ran on3b41fd7c(commit 10): commit 11 changes no kernel. TP=3's ranks ran commit 11's engine source (the first TP=3 run, which found the bug,b7fecb3d's) with two commits our ring of three GB10s needs, which are not part of this pull request (tp3/ring.patchin the receipts): rank 0's NCCL id handed to the other ranks in a file, so that a ring of /32 links bootstraps (#1178 does it upstream's way), and a heat ceiling (TF_HEAT_CEILING: a prompt chunk also waits while its expected rise would pass the ceiling), which changes when chunks run, not what they compute.Environment: GB10s (sm_121, 128 GB), two at TP=2 over a direct RoCE link and three at TP=3 over a ring of direct links, driver 580.178.04, the ranks in Ubuntu 24.04 containers on the efficiency cores (
0-4,10-14) with NCCL 2.30.7, built with Zig 0.17.0 and nvcc 13.3.73 (-Dsm=121), checkpointnvidia/GLM-5.3-Flash-NVFP4@423acf37.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. Decode bybench_openai.py(256 tokens, 3 reps, greedy), cold prompts bybench_cache.py --require-cold(a 32k chat) andreuse_served.py cold(1k to 64k). Seconds, lower is better; a ratio is this head's faster serve over main's slower.fibonacci-raw, s/tokengpu-chat-no-think, s/tokenGreedy decode and the 32k chat gave the same tokens in all five serves. At 16k and 32k main's own two serves differ by 1.8% and 1.1%, which is where the ratios fall below 0.99. The chunk kernels' SASS here is main's loop:
tf_glm_chunks_32has 51 branches, 22 async copies and 20 global loads where main has 50, 22 and 19.Exactness:
tf-cuda-test glm-paged-checkon the checkpoint's first five layers (TF_GLM_LAYERS=5), at TP=2: 7 prompts at the pages' edges (511, 512, 513, 2,047, 2,048, 2,049 and 5,000 tokens, the last reading pooled keys across pages in its sparse pass) and 91 window rows equal their whole-window twins' logits, tokens and drafts byte for byte, with the pages taken out of order (12 jumps) and NaN before use, and the pool idle after; rank 1 made rank 0's pool and mapped its pages through every step. The same on one GPU.TF_GLM_LAYERS=2):glm-ref,glm-synth,glm-lanes-cover,glm-lanes-concurrent,glm-shared-checkandglm-paged-checkpass.--context 131072 --parallel 4, each withbench_concurrent.py --levels 1,2,4 --alone --serial --tokens 256 --reps 1and thenreuse_served.py record --planned --min-prompt 2048:TF_KV_PAGES=all TF_KV_POISON=1 --prompt-cache-gib 4(startup: 3 streams on 504 pages of 512 rows, 180 MiB fixed each, the engine's own sequence idle): 12 cells, every reply equal to its alone run (28/28) and to"draft": false(30/30); the record's 44 requests all cached as planned.TF_KV_PAGES=on TF_KV_POISON=1 --prompt-cache-gib 4(509 pages beside the engine's own sequence): 28/28 and 30/30; 44 requests, all cached as planned. Here a request alone takes the engine's own sequence, so the pages carry the concurrent ones (levels 2 and 4, the record's concurrent group); underallevery request runs on pages.onrecord replayed withTF_KV_PAGESunset and--prompt-cache-gib 0, on this head, whose startup line showsauto's windows at TP=2 ("4 streams at once, 1.77 GiB each at a 131072-token window, of 7.6 GiB left after a 12.2 GiB reserve"), as main's does, and on mainf8d254c5:reuse_served.py compareagainst theonrecord and against theallrecord gives 43/43 replies equal and as cached as planned in all four (the cancelled request is not compared), drafted == plain on the same messages 23/23, and no resumed reply slower than its fresh one. The paged streams answer as main does.TF_KV_PAGES=all TF_KV_POISON=1 --prompt-cache-gib 4 --vision, 305 pages, the repository'sassets/tensorfold-hero.png): a request with drafts, its repeat,"draft": falseand three at once give the same tokens (media_http.py, prompt 2,050), and a chat resumed past its image (6,144 of its 7,605 tokens cached) gives the cold reply's tokens, as do"draft": falseand two at once (media_resume_http.py); both tools are in the receipts'tools/.kv pool idle: promised 0, free 504 + store 0 == total 504(509 underon, 305 with images).tf-cuda-test glm-paged-checkat TP=3 on the same five layers: the same 7 prompts and 91 window rows equal their twins byte for byte, the pages out of order and NaN before use, the pool idle after; ranks 1 and 2 made rank 0's pool and mapped its pages through every step.--context 1048576 --parallel 8,TF_KV_POISON=1on pages:TF_KV_PAGESunset, soautopages at three ranks (startup: "8 streams at once, the engine's own at a 1048576-token window (12.88 GiB, in the model's), 7 more taking 1284 pages of 512 rows (7.99 GiB) as they need them, 131 MiB fixed each, of 24.9 GiB left after a 12.2 GiB reserve"):bench_concurrent.py --levels 1,2,4,8 --alone --serial --tokens 256 --reps 1, 16 cells, every reply equal to its alone run (60/60) and to"draft": false(72/72);reuse_served.py record --planned --min-prompt 2048, 44 requests, all cached as planned; cold prompts of 1k to 64k tokens (reuse_served.py cold, 0.73 s to 43.3 s); and seven 2k prompts sent while a 64k conversation held its rows on pages (63,488 of its 65,536 tokens resumed from the cold run's kept state): all seven equal their alone runs, and none waited for another request to end.TF_KV_PAGES=all(1,237 pages, the engine's own sequence idle, so seven streams at once): 16 cells, every reply equal to its alone run (60/60) and to"draft": false(72/72); in the cells of eight the eighth request waited for a paged sequence (their longest first token 12.0 to 21.2 s, against 1.7 to 2.1 s underauto).TF_KV_PAGES=off --prompt-cache-gib 0 --parallel auto, a window a stream as before: a 1M window leaves room for two ("2 streams at once, 12.88 GiB each at a 1048576-token window, of 24.7 GiB left after a 12.2 GiB reserve"; a named--parallel 8is refused). The record from the serve withTF_KV_PAGESunset, replayed there:reuse_served.py comparegives 43/43 replies equal and as cached as planned, drafted == plain on the same messages 23/23, and no resumed reply slower than its fresh one.TF_KV_PAGES=allwith a 22.8 GiB prompt cache leaves 206 pages (1.28 GiB).kvp_admit.py --fittakes the longest prompt of which one request fits the pool and two do not, 64k at 129 pages each. Three such prompts ran alone, each keeping its state in the prompt cache, then together: all three equal their alone runs, each resumed 63,488 of its 65,536 tokens, and the second and the third started only after the one before ended.kv pool idle: promised 0, free 1284 + store 0 == total 1284(1,237 underall, 206 in the short pool).b7fecb3d, servedallthe same way and found the bug commit 11 fixes: in the first cell of eight, the last request took a slot past the seven paged sequences, and binding it made a 12.88 GiB window on rank 0 before the pages refused it. Rank 0's host fell from 11 GiB free to 4 GiB in two seconds and its memory guard stopped the server, so 53 of that run's 60 alone comparisons failed, none unequal (tp3/found-b7fecb3d/).Registers (nvcc 13.3.73, sm_121, no spills), against main's:
tf_glm_scoresgoes from 129 to 157, as it looks up a page for every 16 pools; at its 128 threads the registers still allow three blocks a SM, as on main, and its shared memory (33,792 bytes) is unchanged.tf_glm_chunks_16goes from 121 to 123 andtf_glm_chunks_32stays at 128;tf_glm_lat_write,tf_glm_pool_keysandtf_glm_dsa_tokensuse two to six fewer.Checks on a GB10 (nvcc 13.3.73,
-Dsm=121), on commit 11: the build,test276/276 steps,test-kv-pool16/16 andtest-cuda-hostpass; the synthetic checkpoint's checks above pass, andglm-paged-checkat TP=1 (7 prompts at the pages' edges and 91 window rows, pages out of order and NaN before use) equals its whole-window twins byte for byte;zig fmt --checkis clean.The kernel contract on a GB10 (nvcc 13.3.73,
-Dsm=121,zig build kernel-contracton3b41fd7c; commits 11 and 12 change no kernel):SUMMARY 148 kernels: 148 pass gates 1-4, 0 fail, 1 known findings; 32 controls: 32 red as planted, 0 missed; 1018 limit problems. Main's on the same GPU:SUMMARY 140 kernels: 140 pass gates 1-4, 0 fail, 0 known findings; 32 controls: 32 red as planted, 0 missed; 1018 limit problems. Gate 5 checks 1,167 registered launches here, main's 1,161 and the six this pull request records: on this GPU each of the six runs with the registers and shared memory of its nvcc 13.3 record (PASS limits live). Every limit problem line is main's, the same in the two runs: 885 SM records main's registered launches lack, and 133 launches whose registers on this sm_121 under nvcc 13.3 are not 13.3's records (23 of them have none). The eight registrations, as the run prints them:tf_glm_lat_writetf_glm_pool_keystf_glm_scorestf_glm_dsa_tokenstf_glm_chunks_16dense, withtf_glm_mergetf_glm_chunks_16sparsetf_glm_chunks_32densetf_glm_chunks_32sparseGate 5's records, read by
tf-cuda-test kernel-contract-limitsfrom fatbins built for sm_80, 86, 89, 120 and 121 by nvcc 13.0.88 and by 13.3.73 (on a tree that registered the six launches without usage): every launch fits its limits on all five. On 13.0 every upstream record with a value is current, so these builds read as upstream's 13.0 records were taken. nvcc 13.2 was not at hand, so a 13.2 build checks these launches against 13.0's records. Registers:tf_glm_lat_writetf_glm_pool_keystf_glm_scorestf_glm_chunks_16tf_glm_chunks_32tf_glm_dsa_tokensThe no-GPU CI's steps under WSL2 on this head (an LF tree from
git archive, a fresh local cache):test-kernel-contract72/72,test-golden, the server and pull security tests (20/20),fuzz-cortheon(14/14), the served Flash Next unittest (8) and the GLM snapshot, KV pool and CUDA host tests pass; of those host tests and ofzig build test(3,463 of 3,479 pass, 15 skipped, commit 11's test among them) onlycuda_cost_identityfails, to WSL2's coarse file times as before (it passes on the GB10). lean: 438 problems, as on main, 0 new;zig fmt --checkon the changed Zig files is clean.Fifteen mutations, each guard broken alone on a copy of commit 11, with
test-kv-pool,test-cuda-hostandtest-glm-snapshots-hostunder WSL2: every one fails a test, and the unchanged copy fails onlycuda_cost_identity. An earlier pass found two guards without a test; commit 10 adds them, and commit 11's test catches the fifteenth.kv_pool.splitwithout the slot-less branch (as before commit 1)kv_pool_test"a family without snapshot slots (snap_bytes 0) ..." andcuda_kv_pool"kv pages: a family without snapshot slots keeps the paged plan's cache ..." (division by zero)kept()without the paged plan'scopied(as before commit 1)cuda_kv_pool"kv pages: a family without snapshot slots ..."autopaging from two rankscuda_native"TF_KV_PAGES: auto pages from three ranks ..."onandalltaken with--kv-cache fp8cuda_native"TF_KV_PAGES with --kv-cache fp8: ..."cuda_kv8_test"GLM FP8 KV: an FP8 cache takes no page table, ..."cuda_pages_test"a paged stream takes its request's pages ... and its end gives them back"cuda_pages_test"a follower maps rank 0's ids, and refuses ..."cuda_pages_test"TF_KV_POISON: every page a stream takes is filled with NaN ..."cuda_pages_test"a follower refuses rank 0's pool when its KV page has other bytes, ..."glm_snapshot_host"a paged sequence's prefix keeps and resumes through its pages in row order, ..."glm_snapshot_host"a paged part's piece from inside a page's span keeps its offset, ..." (commit 10)glm_snapshot_host"while a paged stream is current, a prefix's size is the engine's own sequence's, ..." (commit 10)cuda_pages_test"TF_KV_PAGES=all: a stream past the paged sequences waits for one, ..." (commit 11)Not run, and why:
--drafter): no DFlash2 weights on these hosts, so the paged streams were checked with the MTP head's drafts only.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. (The six kernels pass gates 1 to 4 on a GB10, and their launches have gate 5's records, nvcc 13.0's from every SM's image and 13.3's where it differs;SUMMARYabove.)kv_pool.splitandkept(); the rest cover new code, commit 11's fails with the sequence bound before the pages, and each of the fifteen mutations above fails its test. They run on a fake driver, so they need no GPU.)