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]>
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.
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.
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]>
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).
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.
