HHostingByteSoftware
Ctrl K
ComunidadGuías…

internal/autogen/sizing.md

Rama main · réplica de lectura del repositorio original.

Archivo original ↗Preguntar a KiBot
# 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.