[ZhongUncle]

blob() serves experts from the stage buffers whenever the resident expert arena is opted out of (--mmap-experts) or too small for the expert set. The stage buffers are the pool that assembles each blob from the GGUF's three role slices.

They are allocated with a bare new[], so they are pageable. A cudaMemcpyAsync from pageable memory first copies the data into a driver-internal bounce buffer, on the calling thread, before the DMA can start. Every consumer pays this per expert: the profile prefill's blocking fill (generate.cpp:4114), the lent-slot refills, the adaptive tier's swaps, and the admit path where it runs.

This PR allocates the stage buffers with cudaHostAlloc instead, falling back to pageable per buffer when the driver refuses (pinned memory is scarce on exactly the small-RAM machines this path serves). Nothing else changes: same pool, same claim/reuse policy, same bytes copied — a buffer whose only job is to be the source of an async H2D copy should be pinned.

TL;DR — pinning the stage buffers makes every staged H2D copy cheaper. The cost is ≈0.3 s slower startup (profile prefill ≈7%; ≈0.2 s of it measured pinning cost, the rest unexplained) and ~386 MiB of pinned RAM. It only matters when blob() serves from the stage buffers; otherwise they sit idle. Decode tok/s is roughly flat.

• micro-benchmark, one 1.5 MiB blob: the caller blocks 109–167 µs inside cudaMemcpyAsync from pageable vs 2–3 µs from pinned; end-to-end 219–256 → 144–145 µs (1.5–1.8×), the pinned copy at 10.8 of the card's probed 11.1 GB/s
• engine, --mmap-experts in both arms: the adaptive tier's swap round 4.20–4.30 → 0.90–1.18 ms (3.6–4.8×, four interleaved pairs); the lent-slot refill 165–169 → 104–105 ms for 530 slots (1.6×)
• the profile prefill loses ≈7% (four-run means 3277 → 3055 MB/s): ≈0.21 s of it is the measured pinning of the pool's ~257 buffers, the rest unexplained (Findings)
• scope: on a machine whose RAM holds the whole expert set, the pinned complement (a host mirror of the experts the GPU cache doesn't hold) serves blob() directly and the stage buffers sit idle — this PR is for the configs where they don't

What changes

• claim_stage: new uint8_t[stage_blob_] → cudaHostAlloc(..., cudaHostAllocDefault), with a per-buffer pageable fallback
• the pool's unique_ptr gains a small deleter that frees each buffer the way it was allocated
• two one-time stderr notes make the state observable in the log: the first pinned buffer; the first fallback, with the driver's error string
• cudaHostAllocDefault, not Mapped: a stage buffer is only ever the source of an H2D copy, never addressed by the device
• allocation stays under stage_mu_ exactly where the new[] was; a failed cudaHostAlloc's sticky error is cleared on the spot
• pool growth is concentrated in the first distinct-expert sweep — the profile prefill grows the pool to its ~257-buffer steady state, and afterwards buffers are only recycled, never allocated — so no cudaHostAlloc serialization reaches decode
• the buffers are freed through the same close() path that already cudaFreeHosts the complement and exchange buffers — no new lifetime questions

Measured results

Measured on:

• RTX 3080 20 GB (driver 595.91.07, CUDA 13.2)
• Xeon E5-2680 v4 (no AVX512)
• 62 GB RAM
• Ubuntu 24.04
• Qwen3.8-Flash-Next IQ2_XS pack (48 layers × 512 experts, largest blob 1.51 MB)
• --mmap-experts in both arms (see the scope note), --spec 4 --spec-min-p 0.5 + the MTP drafter

Micro-benchmark — per-copy cost of one 1.5 MiB blob, 64-buffer rotating ring (no buffer is re-read within 64 copies), 512 copies per arm per variant, quiet machine:

source	variant	host µs inside cudaMemcpyAsync	end-to-end µs (GPU events)
pageable new[]	copy-only	166.9	256.3
pinned	copy-only	2.0	145.2
pageable new[]	fill-then-copy	108.7	218.9
pinned	fill-then-copy	3.2	144.4

• fill-then-copy memcpys a fresh blob into the buffer right before its copy — the engine's pattern, where fill_stage assembles the blob and the cache fill copies it right away
• the pageable arm's host time is lower there than in copy-only, likely because the just-written blob is still cache-resident when the driver bounce-copies it
• pinned end-to-end is the copy engine at 10.8 GB/s against the card's probed 11.1 GB/s; pageable end-to-end is the bounce staging plus the DMA, serialized

Engine A/B — four interleaved baseline/patched pairs, 512 generated tokens after an 873-token prompt (raw lines in the appendix):

metric (per pair)	baseline	patched	×
adaptive tier, ms per swap round	4.20 / 4.26 / 4.30 / 4.27	1.18 / 1.17 / 0.90 / 1.17	3.6 / 3.6 / 4.8 / 3.6
lent-slot refill, ms (530 slots)	169.2 / 165.1 / 164.6 / 166.0	104.6 / 104.7 / 104.6 / 104.1	1.62 / 1.58 / 1.57 / 1.59
profile prefill, MB/s (10,244 slots)	3294 / 3262 / 3252 / 3300	3038 / 3040 / 3039 / 3103	0.92 / 0.93 / 0.93 / 0.94
decode, tok/s	36.00 / 40.03 / 39.95 / 40.35	39.75 / 41.01 / 39.11 / 40.83	roughly flat

Findings

1. The swap round's host time drops 3.6–4.8×. The adaptive tier's round loop (generate.cpp:10610) issues one async H2D per swapped expert from blob()'s stage buffer, on its own thread; the logged ms/round times that loop. From pageable memory each issue costs the ~110–170 µs bounce copy measured in the micro-benchmark, from pinned memory ~3 µs. Decode tok/s does not move (finding 4): at this swap cadence (every 4 rounds) the copies are already off the critical path. The gain is headroom on the swap thread, which should matter at a higher cadence or on a slower CPU; the copies are also expected to overlap the MTP drafting that follows (the loop's own comment, generate.cpp:10615).
2. The lent-slot refill gains 1.6×, matching the micro-benchmark's end-to-end ratio. The refill (generate.cpp:10180) queues 530 staged copies and waits once: pageable pays ~117 µs more per slot than pinned (313 → 197 µs), the bounce-staging term.
3. The profile prefill's ≈7% loss is mostly the pool's one-time pinning cost, not a per-copy regression. The fill loop is serial — blob() (three memcpys out of the mapped GGUF) then a blocking cudaMemcpy. The sweep of 10,244 distinct experts grows the stage pool to its ~257-buffer steady state before recycling kicks in; a buffer becomes reusable 256 assemblies after its last use. cudaHostAlloc of 1.5 MiB measures 831 µs here vs 0.2 µs for new[], so pinning the ~257 buffers accounts for ≈0.21 s — about two-thirds of the measured mean slowdown (4.5 → 4.8 s, +0.3 s); the remaining ~0.1 s is unexplained. Per-copy costs are otherwise strictly lower when pinned (findings 1–2), the fill loop is dominated by the mmap read (~440 µs per slot), and the loss is paid once at startup.
4. Decode is roughly flat; the trajectory split in the logs is unresolved. Excluding the cold first run, decode differs by at most ±1 tok/s between arms, within the run-to-run spread. The lookup counts split into groups that mostly follow the arm (the forensics are in the appendix), and the swap totals differ with them, so the patch may have shifted this engine's timing-dependent expert refusals; but the cold baseline run shares the patched group's lookup total, which points to an independent rig state. These runs cannot settle it. Either way, ms/round is normalized per round, so findings 1–2 are unaffected.

Validation

• What was compared: baseline (this PR's parent, 82f46a8) vs patched binaries on four interleaved pairs of identical --mmap-experts runs; the micro-benchmark isolates the mechanism
• Pinning confirmed in the log: every patched run prints FileExpertSource: the stage buffers are pinned (cudaHostAlloc) exactly once; no baseline run does
• The engine's own fill check: slot 0 verified passes in every run (the prefill reads one slot back and compares bytes)
• Full suite: ctest 87/89 — ple_parity and expert_multi_test fail identically on the base commit on this host (the first needs the Q2_0 pack's GGUF, not present here; the second needs AVX512, which this Xeon lacks)

Limits

• Scope: with the default config on this rig the 33 GiB complement is fully pinned and serves every blob() — the stage buffers are idle, and nothing measurable changes. The measured --mmap-experts config stands in for any machine whose arena is smaller than its expert set.
• Pinned footprint: the pool's steady state here is ~257 buffers × 1.5 MiB ≈ 386 MiB of non-pageable RAM — real money on the small-RAM machines this path serves. A driver refusal downgrades only that one buffer to pageable, so under pin pressure the pool runs mixed: slower copies for those buffers, nothing breaks.
• The prefill's one-time pinning cost scales with the pool's growth (~257 buffers here), not with the model. With --prefill off, that growth — and its ~0.2 s — moves out of startup and into the first stage-buffer users: the lent-slot refill alone queues 530 slots, more than the pool's steady state, so the refill would pay it inside its own ~0.1 s window. That configuration was not measured.
• The decode-time admit path (expert_source.cpp:2721) is not measurable in this pack's generate configs: --spec requires the token graph, and the graph decides hits on the device with residency frozen per token — "nothing is admitted during a token" (expert_source.cpp:2391; all four runs logged 0 admitted). That path consumes the same buffers through the same fill_slot, so it stands to gain what findings 1–2 measure; it is cited by code inspection, not measured here.
• One model, one quant, one GPU.

Reproduce

# engine A/B (host with the model; pretokenize via the pack's chat template)
./engine/strata --pack Strata-data/packs/iq2_xs --native <shard1.gguf> --ple-gguf <shard2.gguf> \
  --expert-profile data/expert-profile.bin --expert-cache auto --prefill auto \
  --spec 4 --spec-min-p 0.5 --mtp Strata-data/mtp/rt \
  --max-context 24576 --kv int8 --kv-resident 32768 --mmap-experts \
  --tokens-file PROMPT.tokens --max-new 512 --stats
# the "pre-filled" and "lent slots refilled" startup lines and the "adaptive tier" stat line
# carry the figures; interleave baseline and patched runs

The micro-benchmark is not part of this PR; its parameters and build line are in the appendix, and its ~100-line source can be posted as a comment here if anyone wants to rerun it.

Raw `--stats` lines and the micro-benchmark

Engine A/B, all eight runs — 873-token task prompt, 512 generated tokens, interleaved baseline/patched pairs; both arms --mmap-experts (supports findings 1–4):

baseline r1: pre-filled 10244/10244 in 4.5 s (3294 MB/s) — 530 lent slots in 169.2 ms — adaptive tier 4185 swaps, 4.198 ms/round — decode 36.00 tok/s (14223.8 ms) — hits 293709/326400, 0 admitted
patched  r1: pre-filled 10244/10244 in 4.9 s (3038 MB/s) — 530 lent slots in 104.6 ms — adaptive tier 4078 swaps, 1.177 ms/round — decode 39.75 tok/s (12879.4 ms) — hits 293322/326400, 0 admitted
baseline r2: pre-filled 10244/10244 in 4.5 s (3262 MB/s) — 530 lent slots in 165.1 ms — adaptive tier 3908 swaps, 4.257 ms/round — decode 40.03 tok/s (12791.8 ms) — hits 285173/317760, 0 admitted
patched  r2: pre-filled 10244/10244 in 4.8 s (3040 MB/s) — 530 lent slots in 104.7 ms — adaptive tier 4078 swaps, 1.170 ms/round — decode 41.01 tok/s (12483.5 ms) — hits 293322/326400, 0 admitted
baseline r3: pre-filled 10244/10244 in 4.5 s (3252 MB/s) — 530 lent slots in 164.6 ms — adaptive tier 3908 swaps, 4.297 ms/round — decode 39.95 tok/s (12817.2 ms) — hits 285173/317760, 0 admitted
patched  r3: pre-filled 10244/10244 in 4.8 s (3039 MB/s) — 530 lent slots in 104.6 ms — adaptive tier 3927 swaps, 0.898 ms/round — decode 39.11 tok/s (13092.3 ms) — hits 290363/321600, 0 admitted
baseline r4: pre-filled 10244/10244 in 4.5 s (3300 MB/s) — 530 lent slots in 166.0 ms — adaptive tier 3908 swaps, 4.268 ms/round — decode 40.35 tok/s (12687.9 ms) — hits 285173/317760, 0 admitted
patched  r4: pre-filled 10244/10244 in 4.7 s (3103 MB/s) — 530 lent slots in 104.1 ms — adaptive tier 4078 swaps, 1.171 ms/round — decode 40.83 tok/s (12540.6 ms) — hits 293322/326400, 0 admitted

The trajectory families above are discussed in finding 4: replicas within an arm mostly agree in every logged counter (baseline r2–r4, patched r1/r2/r4); baseline r1's slow decode is the cold first run of the session (14.2 s; every later run 12.5–13.1 s), and its lookup total matches the patched family — so the families' arm correlation is weaker than it first looks.

Micro-benchmark — RTX 3080, 1.5 MiB blobs, 64-buffer rotating ring, 8 reps (512 copies per arm per variant), quiet machine. host µs is wall time inside cudaMemcpyAsync on the calling thread; end-to-end µs is stream time between events bracketing the copy, one synchronize per copy so nothing pipelines (supports the TL;DR and finding 2). Build and run:

nvcc -O2 -arch=sm_86 bench_stage_copy.cu -o bench_stage_copy
./bench_stage_copy          # defaults: 1.5 MiB, ring 64, 8 reps

arm              variant           host us   end-to-end us
pageable new[]   copy-only           166.9        256.3
cudaHostAlloc    copy-only             2.0        145.2
pageable new[]   fill-then-copy      108.7        218.9
cudaHostAlloc    fill-then-copy        3.2        144.4

One-time allocation cost, 1.5 MiB buffers, same rig (supports finding 3): new[] 0.2 µs; cudaHostAlloc 831 µs per buffer when accumulating as the pool does — ×257 buffers ≈ 0.2 s once at startup.

Work with Claude Code

[ZhongUncle]

The micro-benchmark behind the TL;DR's first bullet and finding 2, as offered in the PR body — posted here so the source stays with the PR. It is not part of the diff.

Build and run (any NVIDIA GPU; adjust -arch for yours):

nvcc -O2 -arch=sm_86 bench_stage_copy.cu -o bench_stage_copy
./bench_stage_copy          # defaults: 1.5 MiB, ring 64, 8 reps

// bench_stage_copy.cu - micro-bench for P2: cudaMemcpyAsync H2D from a pageable staging buffer vs a
// cudaHostAlloc'd one, at the IQ2_XS pack's largest expert blob (default 1.5 MiB), over a rotating ring of
// 64 buffers so no single buffer stays L3-resident or branch-predictable (the P1 round-3 lesson).
//
// Two variants per buffer type:
//   copy-only:    the ring is filled once up front; each copy reads a buffer last touched 64 copies ago.
//   fill-then-copy: each buffer is memcpys-filled right before its copy - the engine's pattern, where
//                   fill_stage assembles the blob and the cache fill copies it right away.
//
// Metrics per copy: host wall time inside cudaMemcpyAsync (a pageable source makes the driver stage through
// an internal bounce buffer on the calling thread) and GPU event time on the stream.
//
// Build:  nvcc -O2 -arch=sm_86 bench_stage_copy.cu -o bench_stage_copy
// Run:    ./bench_stage_copy [blob_mib=1.5] [ring=64] [reps=8]

#include <cuda_runtime.h>

#include <chrono>
#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <new>
#include <vector>

static void die(cudaError_t e, const char* what) {
    if (e != cudaSuccess) { std::fprintf(stderr, "%s: %s\n", what, cudaGetErrorString(e)); std::exit(1); }
}

struct Arm {
    const char* name;
    bool pinned;
    std::vector<uint8_t*> bufs;
};

int main(int argc, char** argv) {
    const double mib = argc > 1 ? std::atof(argv[1]) : 1.5;
    const int ring = argc > 2 ? std::atoi(argv[2]) : 64;
    const int reps = argc > 3 ? std::atoi(argv[3]) : 8;
    const size_t n = (size_t) (mib * 1048576.0);

    uint8_t* dst = nullptr;
    die(cudaMalloc(&dst, n), "cudaMalloc");
    cudaStream_t st;
    die(cudaStreamCreate(&st), "cudaStreamCreate");
    cudaEvent_t ev0, ev1;
    die(cudaEventCreate(&ev0), "event");
    die(cudaEventCreate(&ev1), "event");

    uint8_t* fill_src = (uint8_t*) std::malloc(n);
    for (size_t i = 0; i < n; ++i) fill_src[i] = (uint8_t) (i * 2654435761u >> 13);

    Arm arms[2] = {{"pageable new[]", false}, {"cudaHostAlloc", true}};
    for (Arm& a : arms) {
        a.bufs.resize((size_t) ring, nullptr);
        for (int i = 0; i < ring; ++i) {
            if (a.pinned) die(cudaHostAlloc((void**) &a.bufs[(size_t) i], n, cudaHostAllocDefault), "cudaHostAlloc");
            else a.bufs[(size_t) i] = new (std::nothrow) uint8_t[n];
            if (!a.bufs[(size_t) i]) { std::fprintf(stderr, "alloc failed\n"); return 1; }
            std::memset(a.bufs[(size_t) i], 0x5a, n);   // pages resident before any timing
        }
    }

    std::printf("blob %.2f MiB x %d-ring x %d reps (%d copies per arm per variant)\n", mib, ring, reps,
                ring * reps);
    std::printf("%-16s %-14s %10s %10s %10s\n", "arm", "variant", "host us", "gpu us", "GB/s(host)");
    for (int variant = 0; variant < 2; ++variant) {
        for (Arm& a : arms) {
            double host_us = 0.0, gpu_ms = 0.0;
            for (int r = 0; r < reps; ++r)
                for (int i = 0; i < ring; ++i) {
                    uint8_t* src = a.bufs[(size_t) i];
                    if (variant == 1) std::memcpy(src, fill_src, n);   // fill_stage's fresh assembly
                    die(cudaEventRecord(ev0, st), "rec0");
                    const auto t0 = std::chrono::steady_clock::now();
                    die(cudaMemcpyAsync(dst, src, n, cudaMemcpyHostToDevice, st), "memcpy");
                    const auto t1 = std::chrono::steady_clock::now();
                    die(cudaEventRecord(ev1, st), "rec1");
                    die(cudaStreamSynchronize(st), "sync");   // per copy: no cross-copy pipelining
                    host_us += std::chrono::duration<double, std::micro>(t1 - t0).count();
                    float ms = 0.0f;
                    die(cudaEventElapsedTime(&ms, ev0, ev1), "elapsed");
                    gpu_ms += ms;
                }
            const double copies = (double) ring * reps;
            const double us = host_us / copies, gus = (double) gpu_ms * 1000.0 / copies;
            std::printf("%-16s %-14s %10.1f %10.1f %10.2f\n", a.name,
                        variant == 0 ? "copy-only" : "fill-then-copy", us, gus,
                        (double) n / (us * 1e-6) / 1e9);
        }
    }

    for (Arm& a : arms)
        for (uint8_t* p : a.bufs)
            if (a.pinned) (void) cudaFreeHost(p);
            else delete[] p;
    std::free(fill_src);
    cudaFree(dst);
    cudaEventDestroy(ev0);
    cudaEventDestroy(ev1);
    cudaStreamDestroy(st);
    return 0;
}

[Niko1221]

Thanks @ZhongUncle. Ported with your authorship, as an opt-in in 0.1.41: STRATA_STAGE_PIN=1 page-locks the stage buffers (with a 3 GiB available-RAM check unless forced). Decode, 6 interleaved pairs, median tok/s with / without: Tesla P100 with RAM capped at 24 GB and --mmap-experts 13.9 -> 15.0 (+8%, 6/6 pairs); RTX 5070 --mmap-experts +1.2%; RTX 3060 within noise. We had it on by default, but our release check found IQ3_S answers corrupted after a 4,096-token prompt with it on (the output is fine with it off). The likely cause: with page-locked buffers the cache-fill copy really runs asynchronously, so a stage buffer can be reused before its copy has finished; with pageable buffers the driver had copied synchronously and hid it. It stays opt-in until the buffer lifetime is fixed. If you want to look at that part, it is the recycle path of the stage buffers in src/core/expert_source.cpp.

[Niko1221]

Thanks. This is in since 0.1.41 as STRATA_STAGE_PIN=1 (opt-in). It stays opt-in because pinned buffers made the cache-fill copy asynchronous and a stage buffer could be reused before its copy had finished; a fix of the recycle path in src/core/expert_source.cpp would let us make it the default.

[Niko1221]

Thanks for the pull request.

We're resetting the tracker for 0.2.0-beta. There are more than 1,300 open issues and pull requests, too many for us to review properly, so we're closing everything that isn't being worked on right now. Closed here does not mean rejected.

To reopen it takes about two minutes:

1. Rebase on main if it no longer applies, and re-run your checks on 0.2.0-beta.
2. Start the title with its area: [Engine], [GUI], [TUI], [Benchmark], [Docs] or [Other].
3. Fill in the pull request template: the problem and result, and the validation (what you ran, on which revision, what you saw).
4. Comment /reopen. A check runs and either reopens the pull request or lists what is still missing.

The rules are in CONTRIBUTING.md.