proposed_window_mib / proposed_reserve_mib and their two constants go to
research/stage-reserve-proposal/ with a README and their own checks; the
production header keeps reserve_mib / room_bytes, and stage_reserve_test keeps
its real-arithmetic checks unchanged.
Correctness (class E): kept even without a speedup.
Flag B (PR #5096): the graph waits for "flag B >= layer + 1" and flag B is one counter that only rises. A non-empty
plan raises it from a host function on the copy stream, after its uploads. An empty plan (n <= 0) raised it on the
spot, even when the copy stream still held an earlier layer's, or the other token group's, uploads: the counter
jumped past a layer whose bytes had not landed. Verifier::raise_flag_b_ordered() now raises it through the copy
stream whenever copies were queued in this window, and directly only when none were. The same ordering is used at
the four places that publish an empty plan when the pool did not (run, run_slot_rows, batch_serve, service), which
had the same hazard. A failed host-function launch used to be ignored (flag B never rose and the window hung); it
now stops the engine like a failed upload does, after releasing the GPU waits.
Adaptive-copy bookkeeping (PR #4844): a round planned on the pipelined loop whose loop drained before step 3 is
marked resident on the serial rule, which used the lo/hi/bi of an earlier round, so only that range of swaps was
marked and the others were lost (the slot held the new expert, the table still named the old one). The pipelined plan
now names the whole round. The range arithmetic moved into AdaptiveCopyRange (adaptive_range.hpp) so a host test can
walk the rounds.
Tests: tests/core/flag_order_test.cu (delayed host-function copy standing in for a slow DMA, a consumer kernel that
reads staging once flag B allows it, a negative control that shows the old direct raise releases it early, flag only
rises, nothing queued, graph replay twice, every device; passes on both A4000s of z440), tests/program/
adaptive_range_test.cpp (serial batches for 0, 1, many swaps and several caps, the stale-range loss, empty round).
Co-authored-by: AlanJager <[email protected]>
Co-authored-by: uvlad23 <[email protected]>
A pipelined group of a layer split is its own window, so 32 slots on four cards were four groups of 8 and nothing
larger could run. The window limit is now a build option (STRATA_WINDOW_MAX_T = auto | 8 | 16; auto is 16 for HIP,
8 for CUDA; SYCL and Metal never set it): include/strata/kernels/window_limits.hpp.
A window of up to 8 rows runs exactly the kernels it ran before (HIP gfx1201 window 8 against main: 3096 of 3097
kernels identical, the other a debug variant; CUDA sm_86 and sm_60: identical or renamed but for a few prefill kernels
NVVM compiles differently, with identical results). The code for 9 to 16 rows is only reached by a wider window:
- per-row arrays of the verifier, the expert source tables (128 / 256 entries) and the remote / peer expert tables follow
the window; CPU expert jobs stay at 8 tokens (a 9..16-row group is two jobs over the same expert);
- the hyper-connection read is made in slices of 8 rows (its Q8_0 launch would otherwise skip rows);
- the grouped expert kernels run a second pass (SKIP template argument) instead of a 16-entry shared array (51 KB,
over the 48 KB static limit of CUDA);
- resident_plan runs on 128 threads (up to 8 rows, as it was) or 256; MMVQ, bf16 GEMV, shared expert and MoE combine
take 9..16 columns only in a wide build; the sm_60 packed path runs wide calls in passes of 8;
- STRATA_P4_TIMING and STRATA_P4_TRACE (include/strata/program/pipe_probe.hpp): opt-in timing and trace of the
pipelined windows, off unless set.
Tests: src/core/window_limits_test.cpp (target window_limits_test). Real 16-row windows (--batch 64 on four stages,
--batch 16 --batch-groups 1) decode the tokens they decode alone.
0.2.1 made a layer split over exactly two NVIDIA cards with --batch and no --batch-groups run one serial group (measured
then on 2x RTX A4000, IQ3_XXS: +12..16%). With the 0.2.7 code the pipelined groups are faster at 2 and 4 streams on
every box and quant measured, so the default goes back to the pipeline.
Measured (tools/serve_load.py, --batch 4, 256 tokens per request, interleaved pairs, total tok/s, pipelined / serial):
- 2x RTX A6000, IQ3_S, all experts in VRAM: 1 stream 101.7 / 101.7; 2 streams 125.7 / 107.4 (+17%); 4 streams 171.4 / 141.0
(+22%). The v0.2.6 binary on the same box: serial default 104.9 and 137.2, --batch-groups auto 124.6 and 169.6.
- 2x RTX A4000, Q2_0 pack (17037 of 24576 pairs cached, 4 clean pairs): 2 streams 80.2-82.5 / 75.0-75.5 (+7..9%, 4 of 4);
4 streams 110.8-111.5 / 94.8-95.3 (+17%, 4 of 4); 1 stream 63.6-68.7 / 68.2-70.3.
- 2x RTX A4000, IQ3_XXS (13056 of 24576 pairs cached, 3 pairs): 2 streams 73.2-73.7 / 65.8-66.5 (+12%, 3 of 3); 4 streams
100.5-100.7 / 82.9-83.6 (+20%, 3 of 3); 1 stream 58.4-60.5 / 60.7-62.6 (serial about 4% faster, 3 of 3).
The single stream runs on the solo path in both; only the A4000s show the 4%.
What changed: batch_groups.hpp has a pure rule nv_serial_wanted(env value, STRATA_DEFAULTS=026). STRATA_BATCH_GROUPS_NV_SERIAL
set to anything but 0 asks for the serial group, an explicit 0 wins over STRATA_DEFAULTS=026, and with the variable unset or
empty STRATA_DEFAULTS=026 gives v0.2.6's serial group. An explicit --batch-groups (a number or auto) still wins. Three or more
cards, AMD builds and --batch-mtp/--adapt-async handling are unchanged. The serial message names the opt-in.
Tests: batch_groups_test (the rule over the env values and 026, two cards pipeline, the existing serial cases). Real engine on
2x A6000 (IQ3_S), no --batch-groups: the engine says "2 groups of 2 slots" / "2 groups of 1 slot"; with 4 and 2 slots, greedy,
greedy+seeded sampled, slots finishing at different lengths and a cancel in the middle of a draft, every slot's tokens equal its
solo decode and the pipelined build's. A long-prompt slot differs from the pre-merge base build in the same way with the flag
and without it (it is the parent's prompt path, not this change).
Not covered: no run on AMD or on three cards; the 4% single-stream cost on the A4000s is measured on one box.
Problem. #2408 / PR #2409: restoring a large parked conversation on a layer split ended the server with a CUDA out of memory
on a later card. The PR checks the free VRAM of every card against --vram-reserve-mib before a restore. Its own review
points out that a loaded card is always below that reserve, so the check would skip nearly every restore on a split (and
on one GPU). It is not ported.
What a restore allocates. conversation_snapshot_restore writes into the session's existing buffers (pools(st), the GDN
state, the indexer tails); refresh_residency only launches kv_stream_reset or copies into the existing ring. The elastic K/V
growth (kvg_ensure) allocates, but only on one GPU. The one thing a restore can create on a card is the calling thread's
private copy stream for it (conversation_copy.hpp private_stream), which it takes from the pool of finished threads when
there is one (#3350). A failed restore is fatal by contract (the session may be half written). So the need is one stream
and one event, and only for a card whose stream does not exist yet (a snapshot loaded from disk, or a restore on a thread
that never copied on that card).
Change. conv_cache.hpp: kRestoreFirstUseMib (64) and restore_fits(free, stream_ready, n, need): a restore is refused only
for a card that has no copy stream yet and less than 64 MiB free. conversation_copy_stream_ready(device) says whether a
card has one. generate.cpp: before the restore, each card of the run is checked; a snapshot that would not fit goes back to
the cache and the prompt is read again, with a log line. A card that has its stream is never refused, however full.
No sample (a failed cudaMemGetInfo) means the restore goes ahead as it did.
Tests. conv_cache_test: the predicate (cards with a stream are admitted at any free size; the edge at the need; no
sample / no cards / no need). conversation_copy_streams_test: copy_stream_ready before and after the first copy, in the pool
after a thread ends, and the measured cost of a card's first copy stream and event (printed; the test requires under
16 MiB, so 64 MiB leaves a wide margin). conversation_snapshot_test: after the first restore, six more restores of two
images leave the free VRAM where it was.
Not covered. The pure predicate and the generate.cpp code build on z440 (CUDA 13.0); the GPU tests are listed in the stream
notes. No model run: gates/2409-restore-guard.sh (a held-VRAM start on a layer split) shows the skip line and an
unchanged answer on the second turn.
Problem. On a layer split the later cards are left with very little VRAM after loading (#3702: 68 MiB on the last card; #3638:
268 and 666 MiB, then every batched prompt read of more than ~400 tokens ends non-finite after hours; #3797: the weight arena
fails with 320 MiB free and the message does not say who holds the rest). The window-graph eviction (12b0bce95) and the copy
stream pool (#3350) are in already; this is the rest of the code audit of what is allocated after load.
Audit. The table of every lazy allocation after load (what, when, size, how a failure is handled) is in
phaseA-tests/notes-027-gh-eng-d.md. Bugs found and fixed here, each small:
- verify.cpp capture / capture_batch and mtp.cpp finish_capture: the cudaGraphUpload and its sync were only printed. The
upload is a warm-up, but a failure (no VRAM) stayed as the thread's pending error and the next kernel wrapper's
cudaGetLastError ended the engine. The error is now cleared and the first launch uploads the graph.
- verify.cpp capture_commit and mtp.cpp finish_capture: the commit graph called cudaGraphInstantiate directly (no
eviction of older layouts, unlike the window graphs) and reported the status of the end-of-capture call, which is success,
so the message read "no error". capture_commit now goes through instantiate_evicting and reports the instantiate's status.
- instantiate_evicting instantiates into a local and sets the member only on success (what a failed instantiate leaves in
its output is not specified, and a non-null member means "captured").
- prefill.cpp carve: after a failed take of the shared region the sub-allocators had no base and each cudaMalloc'd its
buffer (hundreds of MiB in all) on the card that had just run out. carve now returns at once; the caller reports the
failure as before.
- generate.cpp: a failed set_stage_helper / set_peer is reported and the run goes on, but its CUDA error stayed pending;
cleared.
Found and not changed: relend() re-plans only the first participant (#3638 comment). Read: it runs only from the VRAM
command, which is off on a layer split, so it cannot touch a later stage's loan. Other candidates (ignored cudaEventCreate for
the pipelined fences, unchecked expert copies, cudaFuncSetAttribute flags set before success) are in the notes.
Changes for the report of #3797 and the reserve (#3702).
- include/strata/program/gpu_report.hpp: on a failed weight-arena allocation (first card and each stage) the engine now
also prints, for every card of the run, the free and total memory and the largest other programs that hold memory on it
(NVML, loaded at run time; where the driver does not list them it says so). Text only; the failure is the same.
tests/program/gpu_report_test.cpp pins the text.
- include/strata/program/stage_reserve.hpp: the arithmetic of stage_room (what the split keeps free on a card) moved out of
generate.cpp unchanged; tests/program/stage_reserve_test.cpp compares it with the old expression over a grid. The file
also holds a PROPOSAL for a larger window term on a later card (96 MiB + 6 MiB per further window size + 48 MiB, 174 MiB
for --spec 4) which nothing calls: the default placement is not changed. gates/3702-later-reserve.sh measures it.
- The low-VRAM warning of a later card now adds how many MiB the plan meant to keep free there and how far off that was.
Tests: gpu_report_test, stage_reserve_test (host only, built on z440).
Not covered: no model was run. What the ~700 MiB the plan does not count on the last card is (window graph sets, kernel code
loads) needs the run in gates/3702-later-reserve.sh.
Problem: a layer split with --batch pipelines the slots in groups, and the pipelined groups run no MTP drafts.
--batch-mtp with serial windows leaves all but one card waiting. A multi-stream deployment wants both.
What changed: STRATA_PIPED_MTP=N (with --batch-mtp) keeps the pipeline and drafts in it. A group's window carries
each of its slots' rows: the slot's token and up to N MTP proposals (at most --spec - 1; the first always, each
further one while the draft layer was at least --spec-min-p sure of the one before). The last stage judges each
slot's proposals on their own. The number of rows to keep is known only then, and by then the earlier stages run
other groups' windows, which reuse the scratch the commit replays from. So a small save graph copies those rows
right behind each window (Verifier::batch_launch_rows), and every stage queues the commit of the kept rows after
the verdict (Verifier::batch_commit_rows). The slot drafter catches up on its own rows and drafts the next
proposals as an asynchronous chain (MtpDrafter::chain_launch). Unset, nothing changes.
Changes on top of the PR as sent (review):
- The verifier window size is worked out where the verifiers are built, from the slots that fit and the final
groups, not earlier from --batch as asked. Unset it is the old expression.
- STRATA_PIPED_MTP is decided again once the slots that fit are known. It says it is off, and why, when fewer than
2 groups are left, when STRATA_BATCH_SELF_COMMIT is on, or when a group's window would hold more than
kVerifyMaxT rows. The start line names groups and slots correctly.
- batch_groups::resolve keeps its old mtp_dropped rule; only mtp_piped is new, and it also applies with
--batch-groups given (the PR's fix 1). The "--batch-groups 1" test case now expects "not dropped".
- --remote-expert-opt alone no longer blocks --batch-mtp on a layer split: the object it guards is only made with
a helper expert cache (the PR's fix 2).
- docs/BATCHING.md rewritten with the numbers measured here.
Tests: batch_groups_test (the PR's cases, four "unset is as before" cases); builds on CUDA sm_86 (Linux), SYCL
(the sycl tree has its own verify.cpp and generate.cpp and does not include this code) and MSVC. Model runs:
see docs/BATCHING.md and the gate scripts 4121-*.sh.
Not covered: AMD (HIP) was not built or run; the draft layer on the first stage; helper expert caches.
With --batch-groups above 1 on a layer split the pipelined batch loop never starts a round of the adaptive VRAM tier (the
serial batch windows start one every --adapt-every windows). While the slots are busy the tier only moves when they go back
to the solo path, so a fresh engine keeps its start-up profile for minutes.
STRATA_PIPE_ADAPT_EVERY=N (opt-in; unset, 0 or not a number is off): after N group windows pump starts no new group on the
first stage, lets the groups in flight finish, then runs one round as batch_step does between serial windows, adapt() and
apply_pending(true), so nothing in flight reads a cache slot while experts swap and the new experts have landed before the
next launch. Not while a prompt is read beside the slots (in_prompt_read, its loan of cache slots), not with --adapt-async,
not while the last round's copies are in flight, not with the tier off. The rule is two pure functions in
include/strata/program/batch_groups.hpp (pipe_adapt_every, pipe_adapt_due) with 12 new cases in batch_groups_test. With the
switch set the engine says so once at start and the "strata batch (pipelined ...)" line adds the rounds, the experts
swapped in and the time spent. docs/BATCHING.md documents the switch with the author's table (2x RTX 5090, six streams,
309 to 410 rows/s over eight bursts after a fresh start, two engines per setting).
Reviewed: unset costs one counter increment per group window and one false check where a group starts, and prints no new
line. A due round cannot wait forever: no new group starts, the groups in flight finish, then the round runs. A prompt read
sets in_prompt_read before the reader thread starts (including STRATA_PIPED_CONCURRENT) and clears it after the reader
joined; the drain after the read runs the round if it is due.
Tests: batch_groups_test (the new cases). Not covered and still the author's to show (the PR is a draft): a soak with
prompts read beside the slots, a Linux run, another GPU pair, more than two stages, and exactness settings with the switch
set. The model gate is written down.
--adapt-decay took any number. Above 1 the adaptive tier's usage counts grow without end (inf after a few hundred
adaptations) and the swap gains turn NaN, which breaks the sort that orders the swaps; a negative value flips the counts.
The value is now checked where it is parsed (include/strata/program/adapt_decay.hpp, also in sycl/): a finite value in 0..1
is used exactly as before (0 and 1 included, the range the calibration profile already allows); anything else, NaN and inf
too, runs with the default 0.7 and says so on stderr. A warning instead of the PR's exit 2: a risky pick is reported,
the engine still starts.
adapt_decay_test (25 checks) covers the range, the fallback and finite counts after 100000 adaptations.
Ported from PR #1284 (the check; the refusal became a fallback).
The auto chunk scan only tested the GPU cache. A lent slot's expert comes back from the RAM copy or the model file for the length of a
prompt, and every chunk streams the whole file set; past the page cache that set is re-read per chunk (28-45 GB per 29K prompt on
2x RTX 3060 / 32 GB / IQ2_XS at 8192-token chunks, ~1 GB at 4096). prefill_plan.hpp is the pure model (file set, miss share fitted to
#2824, additive stream + compute + drive time); the serve scan and the generate path refine the VRAM pick with it, after the RAM copy
exists. A set that fits the page cache keeps today's pick. STRATA_PREFILL_HOST_PLAN=0 turns it off; --prefill N is untouched.
On RTX 4060 Ti 16 GB + RTX 3060 12 GB (48 layers, IQ3_S) the first pick was K=2 and the #1616 retry
(every stage at least 3 layers) landed on its own edge, K=3: the 3060 ran 3 layers, 12 GB idle, 452 tok/s of
prefill against 802 at 28 layers (the prompt flows through the stages as a pipeline). For two GPUs the retry
now first asks for a placement in which each stage owns at least floor(layers x its room / all rooms) layers
(split_rules::share_floors, at least 3, trimmed to fit); the three-layer rule stays the second try and the
first pick the last, and the log says which one was taken. Nothing changes where the first pick has no
one- or two-layer stage (equal cards), and three or more GPUs keep the #1616 rule alone.
Reported by @Tworks25 on #2644, with the sweep of layer_split 12..32 (PR #2643). split_rules_test gets the
reporter's rooms (14.2 GiB and 5.50 GiB) and the identical-card and degenerate cases.
Least recent use keeps the newest checkpoints and every turn adds one, so after a few short turns all 6 sit in the
last 1-2K tokens. An agent client that rewrites its history a few thousand tokens back then resumes from the system
prompt: on 2x CMP 50HX with the pi coding agent a 210,528-token prompt that was the same as the session for 204,349
tokens resumed at 9,525 and read 201,003 again (100 s).
With the flag the victim is the checkpoint whose removal leaves the smallest gap for its distance from the newest, so
the kept points thin out geometrically: after 200 turns of 1,000 tokens they sit 0, 1K, 3K, 8K, 43K and 200K before
the end, and an edit 6K back reads 8K again (least recent use: 200K), 42K back 43K. The same number of checkpoints, no
more RAM; never the root, the newest or a pinned one, the tail rule first, the plain policy without a candidate. The
resume is the usual exact-prefix match.
On the box: an 82,593-token conversation grown turn by turn, a turn ~6 steps back rewritten, resumed at 73,610 tokens
in 10.2 s against 40.0 s read in full; with the expert placement fixed (--adapt-swaps 0, --pcie-frac 0.17) the
greedy answer was the same text as the full read. conv_cache_test: the spread cases (a 200-turn simulation against
least recent use, the fallbacks). docs/DETAILS.md: a paragraph after the prompt cache's options.
Squash of SuperDistance's three commits of PR #2050 (decode between the pieces of a read on the pipelined groups, the
piece lines counting the rows the slots decoded meanwhile, the stall watchdog staying on for a read in pieces),
applied after #2326 here: the BYIELD give-way is #2326's give_way lambda, the concurrent block runs first and the
piece path serves a read that STRATA_PIPED_CONCURRENT leaves to it (pictures, configurations where it is off).
Co-authored-by: SuperDistance <[email protected]>
The flag works only through a helper cache (--expert-cache-device1..3), and on a plain layer split it is inert
(MULTI_GPU.md, #1447), so the check looks at the helper caches themselves. #2275's rig runs a plain split with the
flag set, and the mode would have stayed off there for a cache it does not have.
From a review of the concurrent read:
- With external serving batch_poll still finds a window that will not finish (its graph ended or failed with a layer
to ring, or no layer rang for 20 s): the served count is read before the stream query and the ring count after it,
at most one query a millisecond; the last ring's time is an atomic. batch_serve reads a window's row count after its
served count, so a window launched while the doorbell thread was descheduled is not served with the last one's rows.
- #1996's pinned upload rings (0.1.43.1) grew a slot at its first larger upload, a pinned free and allocation on the
read's thread: the drafter's records and the prompt path's routing tables reserve their slots before the read
(PinnedRing::reserve, Prefill::reserve_uploads).
- Off, said once, where serving a doorbell makes a CUDA call (--pcie-mode dma, STRATA_FS_SLOTS, the helper caches) or
the expert source recycles the buffers it hands out (the GGUF in place, unbuffered, STRATA_IO_PREFETCH); a prompt
with pictures is read the usual way. The expert sources' read counters are atomics.
The check is a pure function, concurrent_read_off (batch_read.hpp), with a test in batch_read_test.
- pump's all-idle branch turned the stall watchdog off when the last slot finished mid-read; it stays on while a prompt
is read until the request ends.
- Between windows the doorbell thread slept 50 us, a whole timer tick on Windows (up to 15.6 ms): it now waits on a
condition each launch notifies (STRATA_TEST_BELL_SLEEP_US=N, a test, restores the sleep for an A/B). With nothing
decoding the main thread waits for the read's end or a line instead of sleeping. A pool failure the doorbell thread
set is reported before pump clears the flag for the next window.
- The read's log line ends with the slots' rate and, per stage, its windows and the host time that set each up. The
priority stream's cudaGetLastError runs only under the opt-in.
0.1.41 (0391e478) scaled each stage's miss term by 20 GB/s / its probed link, for 2x 4090 D
with CUDA1 on a physical x1 link (1.6 GB/s: x12, K 22 -> 38, prompt ~350 -> ~1400 tok/s).
On 2x Tesla T10 the same rule moved auto from K=23 to K=14: CUDA1 runs 34 of 48 layers, and as
the prompt path is a pipeline the slower stage sets the rate - 256K prompt 1035 -> 795 tok/s,
decode 46.8 -> 39.5; --layer-split 23 on 0.1.43 restores both (xiake-1 in #2007).
split_rules::link_miss_scale: the scale is the fastest stage's link over this stage's, applied
only from a 4x gap (a riser or chipset x1/x4 slot beside an x16 one), capped at 64. A symmetric
rig keeps the fitted formula whatever its link speed; the x1 case still gets ~x15.
STRATA_SPLIT_LINK_ABS=1 restores the 0.1.41 rule. SYCL has no link term (unchanged).
Test: split_rules_test (the x1 case, equal gen3 cards, a 2x gap, caps, unprobed stages).
Not run on the reporter's T10s: the reply asks for the 'link ... -> miss cost' lines.
On two equal cards every K whose caches hold all the profiled pairs predicts the same window
time, and the search kept the first one tried: the smallest K, which gives the last stage (it
also carries the MTP drafter) the most layers and lets its cache fill the card to the reserve.
2x Arc Pro B70 with IQ3_S: auto K=24, CUDA1 at 31.8 of 31.9 GiB, ~11 tok/s; --layer-split 30
~54 tok/s (agnaylor in #1440). A full xe card does not fail an allocation, it moves buffers to
system memory.
The SYCL search now scores every placement it tries with its smallest spare room (a stage's
room after its session minus every expert of its layers) and, among the placements within
STRATA_SPLIT_HEADROOM_TOL (default 0.2%) of the best predicted time, takes the one with the
most (split_rules::pick_with_headroom). A log line names both picks when it moves the split.
STRATA_SPLIT_HEADROOM=0 keeps the old pick. CUDA/HIP placement is unchanged.
Test: split_rules_test (a 48-layer two-card model of the report: plain K=24 with 664 MiB spare
on the last card, headroom K=25 with 1728 MiB; tolerance and tie rules).
0.1.42 let --batch-mtp run on a layer split and made the --batch-groups auto default give way to it: one serial group. On
2x RTX A4000 (IQ3_XXS, 3 pairs) four streams aggregated 62.2 tok/s against 92.7 on 0.1.41's pipelined groups (-33%, 0/3 pairs
faster), two streams 55.7 against 71.4 (-22%). With --batch-groups 2 (which turns --batch-mtp off) 88.3 and 71.1. The start-up
line says why --batch-mtp is off and how to keep it. --adapt-async 1 still gives one group. Solo decode and the answers are
untouched.
The pipelined slot groups never run the adaptive tier, so --adapt-async 1 with --batch on a layer split and no
--batch-groups resolves auto to one group, like --batch-mtp (batch_groups.hpp, batch_groups_test); --batch-groups G
or auto given on the command line is honoured and turns --adapt-async off through the existing gate.
--batch-groups auto became the default for a layer split with --batch >= 2, but the --batch-mtp gate ran before
the groups were resolved: it saw 1 group, kept batch-MTP on, and the groups were then raised to 2 - pipelined groups
do not run the slots' MTP drafts (#1413). The resolution is now a pure helper (include/strata/program/batch_groups.hpp,
CPU test src/program/batch_groups_test.cpp): with --batch-mtp and no --batch-groups the default is one group and
the log says why; --batch-groups N (N > 1) or --batch-groups auto given on the command line pipelines the groups and
turns --batch-mtp off, as before.
A stage with one layer holds a handful of expert slots (a 12 GB 3060: 512), and the
prompt chunk of every card in the pipeline is clamped to what that cache can lend.
When the pick has such a stage, search again among placements that give every
stage at least three layers; keep the new pick and print both and why. With none
startable the first pick stays. Identical cards pick balanced stages and never
reach the second search.
A prompt read on a layer split with no slot decoding was one run, so the BYIELD of the batch_interleave_test step 4 was
dropped at the end of the read; the long prompt kept its slot and the next BGEN failed with 'ERR BGEN: no such free slot'
(pre-existing since 0.1.40.4). That read now goes in pieces of STRATA_BATCH_YIELD_CHUNKS chunks (default 8, 0 = one run), the
same chunks a single run reads, with the yield check between pieces. batch_read_piece holds the arithmetic; batch_read_test covers it.
On one GPU the prompt path computes the draft layer's K/V for each chunk in batches (Prefill::draft_kv, E-9). A layer
split always ran the drafter's own per-group pass instead, because the drafter sits on the last stage. With
STRATA_SPLIT_MTP_BATCH=1 the last stage's prompt path runs E-9 there. Default off: the default path is unchanged.
The drafter's K/V come from Q8_1 x Q8_0 MMQ instead of its own mmvq, so its drafts can differ; the target's tokens are
decided by the verify window.
Prefill::draft_kv says why it declined (optional out-parameter, also in the SYCL port's prefill.cpp), and the split
logs once what the variable did. Which path runs is a pure header (include/strata/program/draft_kv_plan.hpp) with a
CPU test. tools/split_mtp_batch_parity.py compares the greedy texts of the two arms on a live server. Docs: MULTI_GPU.md.
Builds on #1190. The layer split sized its lend regions (the RAM copies of the experts the prompt path borrows) with
Prefill::bytes_needed before the split's ring override was applied (Prefill::set_ring_override says it must be set
before any bytes_needed), so with STRATA_SPLIT_RING=384 the prompt path borrowed 5396 slots and the copy kept 4662.
The override is now decided first (decide_split_ring); include/strata/program/split_ring.hpp holds the pure helpers
(split_ring_slots, split_lend_slots - the scan's own tail-slot arithmetic, now one function for the region and the
borrow), src/program/split_ring_test.cpp covers them on the CPU. A start line reports borrowed vs kept slots and warns
when kept < borrowed. Docs: MULTI_GPU.md.
The engine's GEN line takes pin=N: the first N prompt tokens are a shared read-only prefix (a long document many
short questions follow). The prompt is read in two parts at N and the checkpoint there is pinned: retention never
evicts it, a parked conversation holding it stays parked, the session file keeps it, and a later request that resumes
from it does not park the branch it leaves, so N queries keep one prefix. One pin at a time. Without pin= nothing
changes.
The server takes "strata_prefix" in a chat, messages or responses request: {"messages": K}, {"message": I,
"chars": N} or {"tokens": N}, resolved to the longest common start of the prompt's ids. A prefix that cannot be
marked is said in the log and ignored; a malformed field is a 400.
Tests: conv_cache_test (pinned eviction), conversation_cache_test (a pinned parked conversation is not evicted),
conversation_file_test (the pinned checkpoint is the one saved), serve.test_server.SharedPrefix.
The retained checkpoints are always a prefix chain - the request loop drops
any checkpoint that is not a prefix of the current prompt, because its
positional cells are gone (the KV cache is one arena, one branch of history
at a time). A radix cache's tree collapsed onto the one branch the session
can hold, and the chain's root is the node a radix cache keeps while its
leaves rotate: in practice the end of the system prompt, which every new
chat of the same client mounts through.
First-in-first-out kept that node only by accident of being first: after
prompt_cache newer checkpoints it was gone, and the next chat re-read the
whole prefix at prefill speed (~30 s per 30K tokens). Now the root is
pinned while the budget allows root and leaf apart (cap >= 2; below that
the old behaviour stands), the rest rotates by least recent use, and a
mount re-dates the checkpoint it mounted through.
The policy is one pure function (conv_cache.hpp) with its scenarios in
conv_cache_test.cpp; checkpoint_save/restore and the state format are
untouched, so STRATA_CKPT_REREAD's token-for-token guarantee holds.
Inference engine, OpenAI/Anthropic API on localhost, optional image input, one-file setup
(START-HERE.bat / setup.sh). Model files are downloaded from Hugging Face by the setup.
Niko1221