# autogen — VRAM sizing & load-plan math How a model's `-ngl` / `--n-cpu-moe` / ctx / KV quant / checkpoint numbers are derived. Read this before touching `plan.go`, `kvcost.go`, `generate_sizing.go` or any `estVramGB` producer. Spawn-time re-sizing lives in [`liveoffload.md`](liveoffload.md). ## Multiple GPUs `gpuset.go`. Every eligible adapter is POOLED into one budget, and the model is split across them with `-sm layer --main-gpu N --tensor-split a,b`. Before issue #4 the sizer picked the largest adapter and left the rest of the box idle: a 12 GB + 16 GB pair budgeted 16. Pooling on its own is a bug, not a feature. With `-sm layer` llama.cpp divides layers **and their KV** by the tensor-split ratio, but the fixed costs (logits/output buffer, CUDA runtime context) live on `--main-gpu` alone. A naive "12 + 16 = 28 GB" therefore OOMs the small card while the sizer reports a comfortable fit. The formulation that holds: ``` splittable_i <= FreeGB_i - fixed_i # per device ratio_i = (FreeGB_i - fixed_i) / sum_j(FreeGB_j - fixed_j) ``` Under that ratio every per-device constraint binds at the same moment, so the pooled budget `sum(FreeGB)` is exactly reachable **provided the fixed costs are charged as overhead**. The existing scalar sizer already does that (`prof.Overhead`), so no vector solver is needed: the main device's whole `Overhead` and one `computeCudaCtxGB` per extra device (`ExtraDeviceOverheadGB`) go in, and `TensorSplit` reads the finished `Overhead` back out. Three details that are easy to get wrong: - **Order.** `prof.TensorSplit` must be computed AFTER the compute buffer is folded into `Overhead`, or the main device is under-charged. - **Device ordinals.** nvidia-smi enumerates by PCI bus id; the CUDA runtime defaults to `CUDA_DEVICE_ORDER=FASTEST_FIRST`. On a mismatched pair those disagree and the split lands on the wrong cards silently, so `emitProfile` emits `env: - "CUDA_DEVICE_ORDER=PCI_BUS_ID"` alongside the flags. - **Eligibility.** An adapter under `minGpuVramGB` (default 3) is dropped. The floor is applied to the device's total **after** the shared system-memory pool is folded in, and only a device that looks integrated (an APU) has that pool counted at all (`settings.sharedMemory`, default `auto`) — a discrete card's host aperture is slower than its own VRAM, so it is never budget. "Looks integrated" is decided in this order: `perf.GpuStat.Integrated`, which is the kernel's KFD verdict (`local_mem_size == 0`) and the only signal that separates a Strix Halo's 16 GB carve-out from a card; a name marker (`ryzen`, `athlon`, `radeon graphics`); then dedicated ≤ 8 GB with shared ≥ 3× dedicated. Without this an APU's 2 GB BIOS carve-out sat under the floor and the only GPU in the machine was invisible to every budget (issue #37). Whether an integrated GPU may be a *split target* beside a real card is a separate, default-off knob (`poolIntegratedGpu`): an APU with no other GPU gets the whole budget regardless. `multiGpu: false` collapses everything (budget, split, guard, wizard seed) back to the single device `MainIndex` picks, which is the pre-issue-#4 behaviour. ## Compute buffer **Modeled, not a flat fudge.** `computeBufferGB` (`generate_sizing.go`) = `runtimeCtxGB()` + `computeGraphGB()`, scaled by `settings.computeBufFactor`. Replaced a flat `ubSoloOh=0.17` that undercounted by >1 GB on large-vocab models and drove VRAM spillover. - **`computeGraphGB` is per-CONTEXT**: logits (`VocabSize*min(ub,computeLogitsTokens)*4`) + activations (`ub*EmbeddingLength*~8*4`). Split out because a second `llama_context` over the same model (the baked-in MTP drafter) allocates its own graph, see the drafter section below. - **`runtimeCtxGB` is per-PROCESS**, charged once per llama-server whatever the context count. Three values, because the three GPU runtimes reserve three different amounts: | Runtime | Constant | GB | Decided by | |---|---|---|---| | CUDA | `computeCudaCtxGB` | 0.3 | GPU vendor (`usingCudaGPU`, NVIDIA in the adapter name) | | Vulkan | `computeVulkanCtxGB` | 0.4 | non-CUDA default | | ROCm / HIP | `computeRocmCtxGB` | 0.8 | backend BINARY path (`usingRocmBackend`) | The non-CUDA figure used to be **0**, on the reasoning that a CUDA number would not transfer. It transfers and then some. **ROCm-vs-Vulkan cannot be read off the GPU**, which is why this is the one sizing input keyed to the executable rather than to telemetry: the same RX 7900 XTX runs either build, and they differ by 0.4 GB. `NoteBackendRuntime` (`vram.go`) classifies `settings.serverExe` by path marker (`rocm`, `hipblas`, a delimited `hip`) and is called from `LoadGenerateFile`, the choke point every sizing path shares, so emit / editor preview / spawn guard never disagree. A path that names no runtime leaves the previous verdict standing — a bare `llama-server.exe` is not evidence of Vulkan — and Vulkan is the non-CUDA default because it is both the installer's pick for a discrete GPU and the cheaper of the two, so an unreadable path can never silently inflate the fleet. Provenance, one PEAK measurement each on the same RX 7900 XTX (dedicated + shared, mid-generation): - **Vulkan 0.4** — b10405-vulkan, Qwen3.6-27B UD-Q4_K_XL, ctx 102400, `-ngl 99`: 20.27 GB dedicated + 2.11 GB shared = 22.38 GB real vs a ~22.4 GB estimate already charging 0.4. Residual ~0. - **ROCm 0.8** — lemonade b1328-llama-windows-rocm-gfx110x, Qwen3.8-27B-GSQ-RCO IQ3_S + baked MTP, ctx 151552, `-ngl 99`: ~0.4 GB *over* a 19.12 GB estimate that already charged 0.4. The mechanism is why this is a model and not a fudge: hipBLASLt/rocBLAS page Tensile kernel code objects for the target arch into VRAM and Vulkan has no equivalent. They load **lazily, on the first real GEMM**, which also explains why the older ROCm numbers that produced the original 0.4 came in low — they were taken on an idle process. The PEAK warning below applies to this constant's own derivation, not just to the graph term. Each side is n=1. A third data point on either runtime should move the number rather than be explained away. - **`computeLogitsTokens` is 1024, not 256.** llama.cpp sizes the output *tensor* by `n_outputs`, but the measured CUDA compute buffer still grows with the physical batch (output projection / cuBLAS workspace tiled over the ubatch). At 256 the estimate was nearly ub-blind (+31 MB going 512→1024), so the sizer never charged the extra ub cost and over-committed — ~0.5 GB spill on Qwen3.6-35B-A3B/8 GB. 1024 (== the common max ub) reproduces the measured +0.5 GB. - **`VocabSize`** comes from `token_embd.weight`/`output.weight` tensor elems ÷ `EmbeddingLength` (`gguf.go`); 0 dims → flat `computeFallbackGB`. - **`DetectGpuCompute` must run before anything sizes a model** — hence its unconditional call in `cmd/quartermaster/quartermaster.go`. It used to sit inside the `-generate` branch, so a serve-only start sized a Vulkan box as CUDA. - `effectiveUb` is shared by sizer and emit, so the charged ub matches the emitted `-ub`. ### Measure at PEAK or you will "prove" it is fat and plan into a spill The model is CUDA-shaped, so scaling `computeBufFactor` down on Vulkan/ROCm looks obviously right — it is a trap. An idle llama-server (loaded, one short prompt) holds ~1.9 GB less than the same process mid-generation: `--ctx-checkpoints` copies (~0.18 GB each, 2 live), the exercised compute buffer, and the MTP draft context only appear under a real prompt. Measured, RX 7900 XTX (b10405-vulkan, Qwen3.6-27B UD-Q4_K_XL, vocab 248320 / embd 5120 / ub 1024 / ctx 102400, `-ngl 99`): **idle** 20.51 GB dedicated (vs 20.37 GB modeled weights+KV+draft — making the 1.10 GB compute term look ~1.6 GB over-charged) but **generating** 20.27 GB dedicated **+ 2.11 GB shared** = 22.38 GB real, against a ~22.4 GB estimate. The estimate was right; the idle comparison was not. This build prints no `compute buffer size` line at any verbosity, so per-process PDH counters (`\GPU Process Memory(*)\Dedicated Usage` **and** `Shared Usage`) are the only measurement — reading dedicated alone hides a spill completely, since demoted bytes leave that counter. ### Windows over-commit degrades silently instead of failing WDDM demotes allocations to system memory when a process exceeds its per-process budget (below physical, and shrinking under pressure from other apps), so the adapter gauge plateaus *under* the cap while the overflow moves to shared — 21.88/23.98 GB dedicated while llama held 2.11 GB shared. Symptom: ~12× prefill collapse (644 t/s at `-ngl 64` → 53 t/s at `-ngl 99`, same model and build) plus desktop stutter as `dwm` gets evicted (0.66 GB shared). Consequences for the sizer: **fitting every layer is not the goal** — 64 layers with ~0.8 GB slack beats 65 layers with none by 12× — and the reachable ceiling is meaningfully below physical, so a plan targeting ~100% of live free VRAM lands on the wrong side of the line whenever desktop usage drifts a few hundred MB. ## Batch flags (`-ub` / `-b` / `--n-cpu-moe`) Bench-validated, not folklore (Qwen3.5-35B-A3B IQ4_XS, 8 GB/32 GB rig): - **`-ub` is prefill-only** (decode flat) and, *on MoE with CPU expert offload*, scales monotonically with **no plateau through 1024** — "smaller ub is faster" is false there. A bigger micro-batch amortises the per-batch PCIe expert fetch (**380 → 647 t/s** going 512 → 1024 at `b=2048`). `computeBufferGB` charges its ~1 GB so the sizer won't over-allocate ctx and spill. - **The dense default is 512, not 1024** (`effectiveUb` keys on `meta.IsMoE`). A fully-GPU-resident dense model has no expert fetch to amortise, so the bigger batch buys nothing: on a 7900 XTX with Qwen3.8-27B-UD-Q4_K_XL at `b=2048`, pp2048 measured **682 / 679 / 667 t/s at ub 512 / 1024 / 2048** — flat-to-inverted, 512 ~2% ahead at every depth (d0 682 vs 667, d16k 543 vs 532, d65k 335 vs 327). Because `computeLogitsTokens` caps the vocab-scaled term at exactly 1024, halving ub halves both that term and activations: **~0.37 GB back** on a 27B/151k-vocab model, better spent as ctx. A fully-resident *MoE* is untested and conservatively keeps 1024. - **`-b` is decoupled from `-ub`, fixed at 2048** (clamped `>=ub`, `<=ctx`). A logical batch above the physical one pipelines more micro-batches per `decode()`, overlapping CPU expert-fetch with GPU compute: **+20–38% prefill, plateau at 2048, zero extra VRAM** (the compute buffer is sized by ub only). Previously `-b == -ub`. - The real decode lever is neither: the emitted MoE placement (`-ngl 99 --n-cpu-moe N`) is **~2× the decode of pure layer-offload** (`-ngl 12`). ## MoE expert share **Two share tables.** `plan.go` has `moeExpertShare` (planner); `generate.go` has `genMoeShare` (generation-side ctx sizing + `forceLowActiveMoE`), which adds `qwen35moe`. Both are only a *fallback*: `effectiveShare` prefers the exact `Metadata.ExpertWeightShare` derived from the tensor section, using the arch table only when that is 0. `plan.go`'s `densePlacement` **intentionally diverges** from the PowerShell original (it fixes a per-layer KV-reserve bug) — see the comment at `plan.go`. ## KV quant **Defaults to `f16`, steps down only under VRAM pressure.** `defaultKvQuant` (`kvcost.go`) is the single default picker used by `emitModel`, `EstimatePlan` and `RenderSoloCmd` (that last one used to hardcode `q8_0` with no MoE branch, so the editor previewed a KV type the config never emitted). - Auto = `f16` for every arch; `q8_0` only when `f16` can't reach `denseMinCtx` inside the model's budget. A quantized KV's cost shows up in long-context recall and multi-turn tool use well before perplexity moves, so a shrunken window is the lesser loss. - `settings.kvQuant` pins one type fleet-wide and skips the decision. - **`bf16` is never the auto pick**: same 2 bytes as `f16` but spends mantissa bits on exponent range K/V activations don't need, and `f16` is the native flash-attention path. - Validation is `ValidKvPair`: llama's `kv_cache_types` (`f32 f16 bf16 q8_0 q5_1 q5_0 q4_1 q4_0`) minus `iq4_nl` (no FA kernel), K and V matched. Anything else — including a typo — resets to the default rather than being emitted and killing the spawn. A per-model/variant override only takes effect if it is itself valid and matched. `GetKvCostModel` gives KV size as `Slope*ctx + Const` GB, SWA/hybrid-SSM aware; `Const` is 0 for plain attention. `GetDenseCtx` is the speed-first dense context picker (avoid offloading just to grow ctx). ## Context checkpoints **Checkpoints reserve VRAM and shrink ctx.** Beyond emitting the `--ctx-checkpoints` count, `checkpointReserveGB` (`generate_sizing.go`) reserves VRAM for them inside `sizeProfile` (split across GPU/CPU KV reserve at partial offload), so a model with checkpoints gets a smaller usable ctx — asserted by `checkpoint_test.go`. Cost is `n * (kvConstGB + perTokGB*step)`: **`kvConstGB` is paid per snapshot**, which is why SWA (gemma) is the expensive case and recurrent/plain (`kvConstGB` ~0) is nearly free. **Count and spacing are tuned together.** Because the const term is per-snapshot but the spacing term is not, rewind coverage is bought with a **wider `--checkpoint-min-step`, not more snapshots**: - `defaultCheckpointMinStep` emits `-cms 1024` for SWA (`swaCheckpointMinStep`), keeps llama's 256 elsewhere; emitted only when it differs from llama's default. - `defaultCtxCheckpoints`: SWA 3 (was 6 — 6 × the whole window state reserved GBs on gemma-class models and crowded out ctx), plain attention 3 at default spacing. - **Recurrent/hybrid emit 2** (`meta.FullAttnInterval > 0` — GatedDeltaNet/SSM, e.g. Qwen3.6). Was `0` (checkpoint restore landed at the wrong position → `non-consecutive token position` spam + full reprocess, llama.cpp #21831) but **fixed in the current build** (genVersion v18, 2026-07-17): verified on qwen3.6-27b — after clobbering the slot with a divergent prompt, returning to a prior 3524-tok prompt restored 3520 tok from a checkpoint (4 reprocessed), zero spam; ckpt=0 fully reprocessed. Hybrid KV is cheap (~1/4 layers full-attn) so 2 is nearly free. - The **emitted flag and the charged reserve share one resolver** (`effectiveCheckpointMinStep`) via `profile.CheckpointMinStep`, so they can't drift. A pinned `Override.CheckpointMinStep` / `VariantSpec.CheckpointMinStep` still wins (variants **inherit** the model-wide spacing, matching the effOv merge at emit — unlike `ctxCheckpoints`, which is standalone per variant). `LiveOffloadArgs` and the config-editor estimate both parse `-cms`. This is the **in-RAM same-process** checkpoint path, separate from the server's slot-cache **disk** path. On hybrids the disk path's *exact* save/restore works (measured 2026-08-19: 32,032 of 32,057 tokens reused across a process restart, prefill 60.6s → 0.35s); only its *partial-prefix seeding* stays gated — see `internal/server/slotcache.md` (`seedSkip`). Repro is `kvcache_probe.py append`, not `swap` (`swap` resends an identical prompt, which needs a rewind and always looks like a miss on these archs). ## Speculative decoding / drafters **DFlash drafters auto-detect and link like MTP sidecars, but DFlash is never auto-selected.** - `discover.go`'s `dflashFileRe` matches "dflash" anywhere in a filename (publishers use it as an infix, unlike MTP's `mtp-` prefix) and pairs it to the same-dir target via `GgufRow.DraftPath`/`DraftKind`/`DraftSizeGB` — same `draftByDir` pass as MTP, different regex + kind. - `effectiveSpec` defaults to the free baked-MTP head chained with model-less ngram (`meta.IsMTP` → `draft-mtp+ngram-mod`, benched better than mtp alone), else plain `ngram-mod`. - A paired DFlash file takes effect **only** via an explicit `spec: draft-dflash` override. It used to auto-default when decode was GPU-bound (+17% on a short flat-prompt MoE bench), but real multi-turn use at the 100k tier craters: DFlash's resident draft weights *and* its own full-context KV compound against a tight budget as the session grows, hitting an oversubscription cliff mtp can't (mtp has no separate weights or KV). Still a legitimate opt-in for short-lived/fresh-context work. - `buildCmdLines` emits `-md`/`-ngld 99` and defaults `--spec-draft-n-max` to **5** for dflash (n-max sweep: ~15% reasoning tg over n=3/4, ties n=6) vs **2** for MTP; `SpecDraftNMax` overrides either. - **Never opt in on dense/CPU-bound models:** on Dense-27B it merely ties mtp on decode, costs 1.43 GB resident, and the draft's own prefill craters pp by −55%. - Any separate draft file is charged via `draftOverheadGB` = real on-disk size + 0.1 GB pad. - A **baked-in MTP nextn layer** (no file) is charged in three parts: its own compute **graph** (`mtpDraftComputeGB`, = `computeGraphGB` at the same ub, since the draft context inherits `params_base` and so projects over the same vocabulary at the same physical batch: on Qwen3.8-27B that logits term alone is 0.47 GiB of a 0.55 GiB buffer), a flat `mtpDraftPadGB` (0.15 GB, now only the second context's non-graph allocations: its scheduler and state buffers), plus a **ctx-scaled KV slope**, `mtpDraftSlopeGB` = `nextn_predict_layers x n_kv_heads x (k_len*bK + v_len*bV)`, carried in `profile.DraftSlopeGB` and added to the per-token slope the sizer solves ctx against. llama-server does not run the nextn head inside the main context: it builds a **second `llama_context` over the same model** at the target's full `n_ctx`, filtered to the nextn layers, so the drafter's KV grows with the window. The old flat 0.34 GB was ~3x too fat at 32k and half the real cost at 160k (Qwen3.8-27B: 1 nextn layer, 0.61 GB of f16 KV at 159744). The graph charge replaced `mtpDraftPadGB` carrying all of it, which was ~4x short: Qwen3.8-27B at ctx 114688 measured 15.82 GiB per-process against 13.95 GiB of modeled components. It rides in `EstimateResult.DraftGB`, not `ComputeBufGB`, and is skipped for a separate draft gguf (already priced by its weights) or when the gguf carries no vocab/embd dims (pad only, as `computeBufferGB` falls back to `computeFallbackGB`). - That second context takes its cache type from `-ctkd`/`-ctvd`, which llama defaults to **f16 regardless of `-ctk`**. `buildCmdLines` therefore emits both for ANY active draft spec (not just when a `-md` file is attached), defaulted to the model's own `-ctk`/`-ctv`, which halves the drafter's KV on a q8_0 model. `Override.KvKDraft/KvVDraft` still win. ## RoPE scaling **The only thing that lifts the trained-ctx ceiling, and it derives its own factor.** Every sizing path goes through `ropeCeiling(meta, ov.RopeScaling, ov.Ctx)` (`rope.go`): `""`/`"none"` clamp to the trained length (a bigger window over untrained positions is garbage, not more context); `linear`/`yarn` raise the ceiling to the requested ctx bounded by `maxRopeFactor` (8×). Four call sites: `generate.go` per model **and** per profile (a variant can enable it alone), `generate_cmd.go` preview, `estimate.go` — plus `EstimateInput.RopeScaling`, so the cogwheel preview and the spawn-time `LiveOffloadArgs` re-sizer size the same window the launch gets. `buildCmdLines` also emits a **derived** `--rope-scale` = `ropeFactor(meta, ctx)` (ratio rounded UP to a half step) whenever a scaling type is set with no explicit factor — a bare `--rope-scaling yarn` leaves llama.cpp on the gguf's own factor (1.0 unless the publisher fine-tuned for extension), i.e. the flag is on and does nothing. An explicit `ropeScale` always wins. ## Pinned custom launch arguments A model's custom launch arguments are constraints on the plan, not only on the argv (`pins.go`; design in [`../ui-svelte/launch-args.md`](../ui-svelte/launch-args.md)). `PinsFromArgs` reads the flags the sizer models (`-c`, `-ngl`/`--n-cpu-moe`, `-ctk`/`-ctv`, `--no-kv-offload`, `--parallel`, `-ub`, `--spec-type`, `--rope-scaling`, `--ctx-checkpoints`, `-cms`), and `emitModel`, `RenderSoloCmdLayers` and `EstimatePlan` fold them in. The emitted occurrence of each flag then says the same thing as the text that replaces it, and `estVramGB`/`estRamGB` describe the launch that actually runs. Two conventions matter: - `-c` is the shared POOL. `profile.Ctx` is the per-slot window, so a pin divides by `--parallel` (`buildCmdLines` multiplies it back out), and `EstimateInput.Parallel` makes the preview charge every slot's share of the KV. - A pinned window is used EXACTLY (`profile.CtxExact`, `EstimateInput.CtxExact`): no `RoundedCtx` floor to a 4096 multiple, the way an auto pick is rounded, and no shrink to the RAM budget in the KV-in-RAM branch (the plan reports the overflow instead). Unmodeled flags (samplers, `-cram`, ...) still win at spawn and do not move the plan. ## Emitted footprint (`estVramGB` / `estRamGB`) **The sizer's footprint is emitted, not just a comment.** Every sizing emitter writes a per-model `estVramGB:` line (`writeEstVram`, `generate_emit.go`) from the same number the estimate already computed — `plan.EstVramGB` for LLMs, the `--max-vram` cap for image, the vllm pool for vllm, size+overhead for embedding/qwentts. That field is the **admission input** for the router's VRAM-budget-aware multi-load (`internal/router/group.go`), paired with the top-level `vramBudgetGB:` (`emitVramBudget`, = `settings.targetVramGB`, withheld when `settings.multiResident: false`). - **CPU-resident classes deliberately emit nothing**: ASR (parakeet), SAM and TTS.cpp run on the CPU, so charging them GPU budget would evict a chat model for something that never wanted VRAM. Absent = unknown, which the router reads as "fall back to the static group policy" — never as 0 GB. - **Consequence: a new emitter that puts weights on the GPU must call `writeEstVram`** or its model can never be co-resident. - `writeEstRam` writes the companion `estRamGB:` (the weights/KV share that does *not* fit on the GPU) — purely informational, surfaced as the Models table's Est RAM column so a partial offload's system-memory cost is visible before loading. Nothing in the router reads it.