mirror of
https://github.com/Heretek-AI/Strata.git
synced 2026-11-08 06:48:28 -10:00
Niko1221
2a9a67c174
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.