Niko1221/Strata
1
0 Fork 0
mirror of https://github.com/Niko1221/Strata.git synced 2026-11-07 18:35:04 -10:00
Code Issues Packages Projects Releases 1 Wiki Activity

Resident RAM mode with a layer split

Browse Source 
--resident-experts (the low-RAM mode's resident variant) ran as plain
--mmap-experts as soon as a layer split was set. It now keeps its RAM copy on
a split:

- the experts every stage's cache holds are left out of the RAM copy, not only
  CUDA0's;
- when the rest does not fit whole (#467), the copy keeps the hottest by the
  whole expert profile. Each stage only keeps its own layers' share of the
  profile, and ranked by CUDA0's share alone the copy held layers 0..K-1's
  experts and read every later layer's from the files (4.0 GiB in RAM instead
  of 21.2 on the PC below);
- an adaptive swap copies the evicted expert back into RAM from the cache of
  the card that owns its layer, on that card's stream. host_res holds slot
  numbers of whichever cache owns the layer: reading a later stage's slot out
  of CUDA0's cache put another expert's bytes in RAM, and after the first
  swaps the answers turned to garbage. (Latent until now: the guard kept the
  mode off a split.)
- the decode profile prints every stage's GPU profile, not only CUDA0's.

Remote expert caches (--expert-cache-remote) are still refused with it.

Swift 1.5 IQ3_XXS at 160K, q4_0 KV, RTX 4060 Ti (layers 0-19) + RTX 5080
(20-47), i9-14900KF, 32 GB of RAM, Windows 11, four greedy prompts, decode
tok/s (first round / fourth):
- 5080 alone, --resident-experts: 25 / 29
- split, --mmap-experts (what --resident-experts became): 32 / 64
- split, --resident-experts (22 GiB of experts locked in RAM): 71 / 69
With --pcie-frac 0 --adapt-every 0 the split's greedy output is the same as
--mmap-experts' (4/4 prompts). With adaptive swaps on, 33 varied requests at
temperature 0.6 (thinking on and off, five languages, code, each checked for
degenerate output) exchanged 81,783 experts with the two caches and gave no
degenerate answer.
This commit is contained in:
Francesco Albano authored and Niko1221 committed 2026-10-05 11:03:50 +02:00
1 parent 5c4105c073
commit fe9c10cad7
2 files changed
+92 -29
Download Patch File Download Diff File Expand all files Collapse all files

No files matched your search

docs/MULTI_GPU.md
+17
View File
@@ -102,6 +102,23 @@ it goes).
A card holding more experts can change which experts run on the GPU, so the output can differ slightly from a run
without it.
**The resident RAM mode works on a split** (`--resident-experts`, the low-RAM mode). The RAM copy of the experts
leaves out the ones every card's cache holds, not only the first card's, and when the rest does not fit whole it keeps
the hottest by the expert profile over all the layers. An adaptive swap copies the evicted expert back into RAM from
the card that owns its layer. Before, `--resident-experts` with a split ran as `--mmap-experts`. Swift 1.5 IQ3_XXS at
160K (q4_0 KV, `--prefill 4096`, `--spec 4` with the stock draft layer), RTX 4060 Ti (layers 0-19) + RTX 5080 (20-47),
i9-14900KF, 32 GB of RAM, Windows 11, four greedy prompts at a time, decode tok/s:
| | first four prompts | after three more rounds |
|---|---|---|
| 5080 alone, `--resident-experts` | 25 | 29 |
| split, `--mmap-experts` (what `--resident-experts` became on a split) | 32 | 64 |
| split, `--resident-experts` (22 GiB of experts locked in RAM) | 71 | 69 |
The split with `--mmap-experts` catches up once the OS file cache holds the experts, on a PC with nothing else
running; the resident copy is there from the first request and stays locked when other programs need the RAM. With
`--pcie-frac 0 --adapt-every 0` the split's greedy output is the same with either mode.
**A separate VRAM reserve for the later cards:** `--vram-reserve-later-mib N` (default: `--vram-reserve-mib`'s value).
The card that drives the monitors needs more headroom than one that drives none; with the display on the last card,
`--vram-reserve-mib 300 --vram-reserve-later-mib 1800` gives the first card's cache that VRAM.
src/program/generate.cpp
+75 -29
View File
@@ -209,13 +209,23 @@ using Clock = std::chrono::steady_clock;
// reads the file. Swaps that need no exchange (`out` in the lend region is held in RAM already; or `in` is not) go
// on as before; ones beyond the buffers' room wait for a later round. Runs on the adaptive tier's thread while the
// GPU commits and drafts: the copies back are on its stream, and waited for before the refills are queued.
template <class Swap>
bool resident_stage_swaps(strata::core::FileExpertSource& src, strata::core::ExpertCache& cache,
const std::vector<int32_t>& host_res, int64_t n_expert, std::vector<Swap>& swaps,
cudaStream_t stream) {
//
// With a layer split, `host_res` holds slot numbers of WHICHEVER cache owns the layer, so the copy back must read the
// owning card's cache on that card's stream: `locate(layer)` names them (cache, stream, device; -1 = CUDA0). Reading
// a later stage's slot number out of CUDA0's cache copied another expert's bytes into RAM - the resident mode with a
// split then computed some experts from the wrong weights after the first adaptive swaps (garbage output).
struct SwapHome {
strata::core::ExpertCache* cache;
cudaStream_t stream;
int dev;
};
template <class Swap, class Locate>
bool resident_stage_swaps(strata::core::FileExpertSource& src, const std::vector<int32_t>& host_res,
int64_t n_expert, std::vector<Swap>& swaps, Locate locate) {
if (!src.complement_ready() || swaps.empty()) return true;
struct Staged { int32_t layer, in, out; int64_t q; };
std::vector<Staged> staged;
std::vector<SwapHome> used;
std::vector<Swap> kept;
kept.reserve(swaps.size());
for (const Swap& s : swaps) {
@@ -224,17 +234,35 @@ bool resident_stage_swaps(strata::core::FileExpertSource& src, strata::core::Exp
if (q >= src.exchange_capacity()) continue;
const int32_t slot = host_res[(size_t) s.layer * (size_t) n_expert + (size_t) s.out];
if (slot < 0) continue;
if (cudaMemcpyAsync(src.exchange_buffer(q), cache.device_slot(slot),
const SwapHome home = locate(s.layer);
const strata::core::OnDevice on(home.dev);
if (const cudaError_t e = cudaMemcpyAsync(src.exchange_buffer(q), home.cache->device_slot(slot),
(size_t) strata::kernels::cpu::expert_layout().blob_bytes(s.layer), cudaMemcpyDeviceToHost,
stream) != cudaSuccess)
home.stream); e != cudaSuccess) {
std::fprintf(stderr, "strata: copy back of layer %d slot %d (device %d) failed: %s\n", (int) s.layer,
(int) slot, home.dev, cudaGetErrorString(e));
return false;
}
bool seen = false;
for (const SwapHome& u : used) seen = seen || u.stream == home.stream;
if (!seen) used.push_back(home);
staged.push_back({s.layer, s.in, s.out, q});
kept.push_back(s);
}
if (!staged.empty()) {
if (cudaStreamSynchronize(stream) != cudaSuccess) return false;
for (const SwapHome& u : used) {
const strata::core::OnDevice on(u.dev);
if (const cudaError_t e = cudaStreamSynchronize(u.stream); e != cudaSuccess) {
std::fprintf(stderr, "strata: copy back sync (device %d) failed: %s\n", u.dev, cudaGetErrorString(e));
return false;
}
}
for (const Staged& x : staged)
if (!src.stage_exchange(x.layer, x.in, x.out, x.q)) return false;
if (!src.stage_exchange(x.layer, x.in, x.out, x.q)) {
std::fprintf(stderr, "strata: stage_exchange refused layer %d in %d out %d\n", (int) x.layer,
(int) x.in, (int) x.out);
return false;
}
}
swaps.swap(kept);
return true;
@@ -1612,19 +1640,10 @@ int main(int argc, char** argv) {
}
const bool remote_caches = o.expert_cache_remote[0] > 0 || o.expert_cache_remote[1] > 0 ||
o.expert_cache_remote[2] > 0;
if (o.resident_cpu_experts && !o.layer_split.empty() && !remote_caches && o.resident_soft &&
!o.resident_cpu_explicit && o.resident_budget == 0) {
// #364 #384: setup's --resident-experts with a layer split (--gpus at start, or a config edited by hand) runs
// as the plain mmap mode - the placement those users measured 1.3-1.6x faster than one GPU - instead of
// refusing. Exactly --mmap-experts: nothing else reads these flags (the headroom only sizes the copy).
std::fprintf(stderr, "strata generate: WARNING: the resident RAM mode (--resident-experts) does not support a "
"layer split yet: the experts the GPUs do not hold are read through the OS file cache "
"(--mmap-experts), and RAM may fill up during long prompts\n");
o.resident_cpu_experts = o.resident_pin = o.resident_soft = false;
o.resident_headroom = 8ull << 30;
}
if (o.resident_cpu_experts && (!o.layer_split.empty() || remote_caches)) {
std::fprintf(stderr, "strata generate: --resident-cpu-experts does not support layer splits or remote expert caches\n");
// A layer split keeps the resident RAM mode: every stage's GPU cache is left out of the RAM copy, and an adaptive
// swap copies an evicted expert back from the card that owns its layer (resident_stage_swaps).
if (o.resident_cpu_experts && remote_caches) {
std::fprintf(stderr, "strata generate: --resident-cpu-experts does not support remote expert caches\n");
return 2;
}
// the helper-GPU expert caches (--expert-cache-remote, docs/SECOND_GPU.md): CUDA1..3 on one GPU; with a layer
@@ -2857,7 +2876,12 @@ int main(int argc, char** argv) {
while (st < (int) split_at.size() && l >= split_at[(size_t) st]) ++st;
return st;
};
// the whole ranking, hottest first, before each stage takes its own layers' share: the resident RAM copy ranks
// the experts no GPU holds over ALL the layers (with CUDA0's share alone, a copy that did not fit kept only layers
// 0..K-1's experts and read every later layer's from the files)
std::vector<std::pair<int32_t, int32_t>> profile_all;
if (multi_gpu) {
profile_all = profile;
std::vector<std::pair<int32_t, int32_t>> mine;
for (const auto& pr : profile) {
const int st = stage_of(pr.first);
@@ -4551,16 +4575,24 @@ int main(int argc, char** argv) {
const int64_t k = plan_lend(chunk);
if (k > 0) lend_from = xcache.slots() - k;
}
bool resident_ok = src.pin_cache_complement(xcache, err, o.resident_pin, {}, lend_from, o.resident_headroom,
o.resident_budget, &profile);
// a layer split: the experts every later stage's cache holds are left out of the RAM copy, as CUDA0's are
// (with them the copy keeps no lend region: pin_cache_complement turns the loan off)
std::vector<std::pair<int32_t, int32_t>> stage_pairs;
for (auto& st : stages)
for (int64_t l = st->lb; l < st->le; ++l)
for (int64_t e = 0; e < g.n_expert; ++e)
if (st->cache.slot_of(l, e) >= 0) stage_pairs.emplace_back((int32_t) l, (int32_t) e);
const std::vector<std::pair<int32_t, int32_t>>& rank_all = profile_all.empty() ? profile : profile_all;
bool resident_ok = src.pin_cache_complement(xcache, err, o.resident_pin, stage_pairs, lend_from,
o.resident_headroom, o.resident_budget, &rank_all);
std::string whole_err;
if (!resident_ok && o.resident_soft) {
// #467: the whole complement does not fit - keep what does, the hottest by the profile, through the #403
// budget path (sized by the RAM alone) instead of none: the misses outside it read the same file bytes
// the mmap fallback reads, so the answers are unchanged. Nothing pinned: the old fallback below.
whole_err = err;
resident_ok = src.pin_cache_complement(xcache, err, o.resident_pin, {}, -1, o.resident_headroom,
strata::core::FileExpertSource::kResidentWhatFits, &profile);
resident_ok = src.pin_cache_complement(xcache, err, o.resident_pin, stage_pairs, -1, o.resident_headroom,
strata::core::FileExpertSource::kResidentWhatFits, &rank_all);
if (resident_ok)
std::fprintf(stderr, "strata generate: WARNING: the whole resident RAM mode does not fit (%s); %.2f "
"GiB of the experts the GPU does not hold, the hottest by the expert profile, are "
@@ -5531,7 +5563,16 @@ int main(int argc, char** argv) {
}
std::sort(swaps.begin(), swaps.end(), [](const Swap& a, const Swap& b) { return a.gain > b.gain; });
if ((int) swaps.size() > o.adapt_swaps) swaps.resize((size_t) o.adapt_swaps);
if (!resident_stage_swaps(src, xcache, host_res, g.n_expert, swaps, adapt_stream)) return false;
// the copies back read the cache of the card that owns each layer (see resident_stage_swaps)
auto home_of = [&](int64_t layer) -> SwapHome {
const int stn = multi_gpu ? stage_of(layer) : 0;
if (stn > 0) {
GpuStage* gs = stages[(size_t) stn - 1].get();
return {&gs->cache, gs->adapt_stream, gs->dev};
}
return {&xcache, adapt_stream, -1};
};
if (!resident_stage_swaps(src, host_res, g.n_expert, swaps, home_of)) return false;
if (pin_blobs_on()) { // lock the batch's source pages (a file-backed arena on AMD; see pin_blobs)
std::vector<std::pair<uintptr_t, uintptr_t>> spans;
for (const Swap& s : swaps)
@@ -7224,8 +7265,11 @@ int main(int argc, char** argv) {
(d1.actq - ds0.actq) / w, (d1.jobs - ds0.jobs) / w, (d1.run - ds0.run) / w,
(d1.host - ds0.host) / w, dt_commit / w, dt_draft / w, (d1.misses - ds0.misses) / (w * L),
(d1.entries - ds0.entries) / (w * L), (d1.hits - ds0.hits) / (w * L), (d1.pcie - ds0.pcie) / (w * L));
const std::string pr = ver.profile_report();
if (!pr.empty()) std::fprintf(stderr, "strata decode GPU stages (ms/window):%s\n", pr.c_str());
for (int st = 0; st < n_stages; ++st) { // every stage's GPU profile, not only the first card's
const std::string pr = stage_ver(st).profile_report();
if (!pr.empty())
std::fprintf(stderr, "strata decode GPU stages, stage %d (ms/window):%s\n", st, pr.c_str());
}
}
if (!cancelled) {
// a prompt stopped halfway leaves the session somewhere between two chunks: nothing to continue from
@@ -7949,7 +7993,9 @@ int main(int argc, char** argv) {
}
std::sort(swaps.begin(), swaps.end(), [](const Swap& a, const Swap& b) { return a.gain > b.gain; });
if ((int) swaps.size() > o.adapt_swaps) swaps.resize((size_t) o.adapt_swaps);
if (!resident_stage_swaps(src, xcache, host_res, g.n_expert, swaps, adapt_stream)) {
// this (non-serve) path refills CUDA0's cache only, so its copies back read CUDA0's cache too
if (!resident_stage_swaps(src, host_res, g.n_expert, swaps,
[&](int64_t) -> SwapHome { return {&xcache, adapt_stream, -1}; })) {
std::fprintf(stderr, "strata generate: an adaptive refill failed (copying evicted experts back)\n");
return false;
}