vllm.cpp reads a number of environment variables. Most are kernel-internal
tuning and bisect knobs that belong to the engineering record, not to a
deployment; this page documents the subset that changes user-visible behavior or
is a supported operational lever. Everything else is enumerated on the
kernel-internal allowlist (scripts/env-doc-allowlist.txt) and kept honest by a
CI check (see Keeping this reference honest).
Unless stated otherwise, a flag-style knob is read as on when its value is a
non-empty, non-0/false/off string, and the listed default applies when the
variable is unset. None of these are required to run the engine.
The weight-residency knobs VT_GGUF_MMAP, VT_GGUF_PREFAULT,
VT_MOE_EXPERT_STREAM, VT_MOE_EXPERT_STREAM_SLOTS,
VT_MOE_EXPERT_STREAM_SLOT_BYTES and VT_DEVICE_WEIGHT_BUDGET_BYTES, and the
four placement knobs VT_PLACEMENT_OVERRIDES, VT_CPU_MOE, VT_N_CPU_MOE and
VT_PLACEMENT_FIT, can also be
set in the server's JSON config, under the vllm_cpp key of --offload-config.
That is the documented deployment surface; these variables are the
override:
environment variable > --offload-config's vllm_cpp key > built-in default
VT_CPU_MOE, VT_N_CPU_MOE, VT_PLACEMENT_OVERRIDES and VT_PLACEMENT_FIT
are ENG-HYBRID-PLACEMENT (#149).
They are read, validated, desugared, resolved to a per-layer decision, and acted
on: a placed layer runs its routed experts on the placement device, with the
hidden state crossing at the layer boundary.
Three limits, and each of them bites.
Five architectures, not all of them. Qwen3MoeForCausalLM, Qwen3.5/3.6,
Nemotron-H, DeepSeek-V2 and Kimi-Linear route through one shared seam and honour
these variables. The rest do not, for reasons that differ and are worth telling
apart: Laguna runs its expert GEMMs on the device but presents a per-token
host-float FFN boundary, so it has no [T,H] block to hand the seam;
Gemma4's expert path accumulates into a caller's buffer rather than returning
one; DeepSeek-V4 runs its experts on the host, where a placement has nothing
to move; and GLM-5-Next, dots3-note, Kimi-K3 and qwen4_exp have no reachable MoE
forward yet. On all of those the variables are silently inert.
The TARGET is the CPU. The engine may run on any device; the destination may not. A placement naming an accelerator is refused at startup, because the queue is created once for the process and a non-CPU queue would leak its stream.
No model has been run with a placement in force. The mechanism is gated by unit and round-trip tests — the round trip is byte-identical to computing in place, and mutation-proven — and the end-to-end token gate is owed. Treat a resolved placement on the install line as what was asked for, not as a measured result.
docs/FEATURES.md carries the capability as partial for these reasons and is the
row to trust for scope.
The polarity is deliberate. These variables exist so a benchmark arm is
switchable without restarting the server with a new document, so an override that
could not turn a configured knob back off would be useless. VT_X=0 therefore
beats a config true, and the engine prints one line at startup naming the
document it installed and a second naming every variable that would WIN over a
field of it — a configuration silently overridden by something exported weeks ago
is the one way this precedence hurts. "Would win" is exact for the three boolean
knobs, where any value wins. For the two COUNTS it is narrower than presence: an
empty, garbage or non-positive value falls through to the config under the
tolerant parse those readers have always used, so VT_MOE_EXPERT_STREAM_SLOTS=banana
overrides nothing and is not reported as an override.
VT_DEVICE_WEIGHT_BUDGET_BYTES is narrower again, and in a different direction.
Its grammar is decimal digits only — no sign, no space, no trailing text — so a
value outside it falls through to
vllm_cpp.device_fit.weight_budget_bytes and then to the device probe. Inside
that grammar every value wins, 0 included: 0 is this knob's suppression
spelling rather than an absent one, so it is an override and is reported as one.
It joined the config surface after the other five
(#1127).
See the expert streaming guide for the
config form and
.agents/specs/weight-residency-config.md
for why the extension is a namespaced key rather than a field on vLLM's mirrored
OffloadConfig.
VT_MOE_EXPERT_STREAM_STATS_EVERY is deliberately not a config key. Every
knob above changes what memory the process reserves or where a weight lives,
which is a deployment decision; that one changes only how often a diagnostic line
reaches stderr. It is the instrument, not the configuration.
These change how the engine runs and have no CLI flag (or complement one).
| Variable | Default | What it does |
|---|---|---|
VLLM_CPP_CPU_THREADS |
hardware concurrency | Overrides the CPU threadpool width. The single most useful CPU-deployment knob; there is no CLI flag for it |
VT_CPU_SPIN_ROUNDS |
4096 on aarch64, 256 elsewhere |
How many relax rounds a CPU-threadpool waiter spins before yielding its core, in Threadpool::Barrier and Threadpool::PollForWork. A waiter that never yields costs a full scheduler timeslice per dispatch as soon as the pool is wider than the cores available to it, and a stock run reaches that because the pool defaults to hardware concurrency while the process has other runnable threads. 0 restores the never-yield spin for a same-binary A/B. Neither setting changes any computed value: the yield is a scheduling hint only |
VLLM_PREFIX_CACHING_HASH_SEED |
0 (fixed) |
Seed for the prefix-cache block hash, mirroring vLLM's PYTHONHASHSEED. random makes block hashes non-deterministic across processes, which takes any persisted or shared KV cache to a 0% hit rate. Keep it fixed if you rely on cross-process prefix reuse |
VLLM_KV_EVENTS_USE_INT_BLOCK_HASHES |
1 (on) |
Whether published KV-cache events carry block hashes as an int (the low 64 bits of the sha256 digest) rather than the raw 32 bytes, mirroring vLLM's env of the same name and its default. Set 0 to publish the raw bytes. Only affects the KV-cache event payload (--kv-events-config); it does not change the internal block hashing or the cache itself |
VT_TT_KEEPQUANT_FUSED |
1 (on) |
Tenstorrent backend: decode keep-quant words through the fused whole-extent kernel (one program per chunk launch, common runtime args) instead of the ~20-program eltwise chain. 0 restores the chunk chain; the arm also declines by name for over-L1-budget planes and non-golden-gated encodings. The fused path is bit-exact to the chain and cuts the captured trace ~6-90x per launch |
VT_TT_KEEPQUANT_MM_CHAIN |
0 (off) |
Tenstorrent backend: force the MatmulBT operand-prep chain fallback (1), bypassing the wave-2 fused launch even where it is golden-gated. A/B rollback for the fused matmul path; the chain is bit-exact to the fused path |
VT_TT_KEEPQUANT_INT8DOT |
1 (on) |
Tenstorrent backend: serve the f32-out dense keep-quant arm with the grouped int8-dot decode kernel (one captured launch per matmul, activation quantized once on-core) for the encodings outside the unconditional set, instead of the per-chunk E=1 grouped dequant chain. 0 restores the W4a grouped f32-exact arm (the same-binary rollback; outputs differ — the int8-dot domain is op-level bit-exact to vt::cpu::BlockVecDot, not f32-exact to the chain). Default flipped 2026-10-04 on the live 27B band evidence: both arms 16/16 in-band at 500 mnats, TPOT 4.69x (docs/bench-evidence/tt-int8dot-band-16p-20261004.md) |
VLLM_CPP_REGION_CAPTURE |
0 (off) |
Tenstorrent backend: capture the 27B decode graph as one region per layer instead of one whole graph (1 enables). Each region is bounded to 50 MiB of trace staging and an over-budget region declines to eager by name; the whole-graph arm is still attempted first when the fit predicate passes. The default path is unchanged — the arm is insurance for the whole-graph trace-fit wall and for devices where whole-graph capture cannot fit |
VLLM_PLUGINS |
unset (load all registered) | Comma-separated allowlist of general plugins to load in LoadGeneralPlugins(), mirroring vLLM's VLLM_PLUGINS. Unset loads every registered plugin; an empty string loads none; a list loads only the named plugins. A plugin that throws is logged and skipped (the load never aborts the engine). See .agents/specs/plugin-system.md |
VT_LMCACHE_HOST |
127.0.0.1 |
Default LMCache server host for the lm:// connector. The kv_connector_extra_config.host key overrides it. See KV offload |
VT_LMCACHE_PORT |
65432 |
Default LMCache server port. The kv_connector_extra_config.port key overrides it |
VT_LMCACHE_HASH_ALGO |
blake3 |
Default LMCache key-derivation algorithm. Set vllm (alias sha256_cbor) for byte-for-byte interop with a real vLLM + LMCache peer. The kv_connector_extra_config.hash_algo key overrides it |
VT_SERVER_MAX_PROMPT_CHARS |
unset (no fixed cap) | An optional operator ceiling, in bytes, on the rendered prompt for /v1/chat/completions. The server refuses a larger prompt with HTTP 400 BadRequestError and a message that names this variable. Unset or 0 sets no fixed cap. Without it, the server still refuses a chat prompt that cannot fit the context: before tokenization, when the rendered prompt, tools and template text included, is longer than max_model_len times the longest token in the vocabulary, in bytes; and after tokenization, when it is longer than max_model_len tokens. No prompt that fits in max_model_len tokens is longer than that byte bound, with one exception: a SentencePiece tokenizer that sets fuse_unk and has no byte fallback encodes a run of unknown characters of any length as one <unk> token, so a prompt made of such runs can fit in max_model_len tokens and still be refused by the byte bound. vLLM's max_chars_per_token bound has the same exception. A value you set here can refuse a prompt that fits the context |
VT_SERVER_SSE_PING_S |
0 (off) |
Seconds between SSE comment keepalives (:\n\n) on silent streams; <=0 disables. OFF is the default because vLLM's streaming endpoints emit no comment frame at all, and some SSE clients — including vLLM's own vllm bench serve — cannot resynchronise after one, silently counting the request FAILED while the server completes it. Set a positive value (clamped to 600) only behind a proxy whose inactivity timeout would otherwise drop a long prefill, and expect strict clients to break. See #931 |
VT_SERVER_MAX_NEW_TOKENS |
4096 |
Caps a chat request's max_tokens value. Set 0 to disable the cap |
VT_TT_HOST_FREE_DECODE |
unset (host-free eager decode) | Tenstorrent-only decode mode: residual-RMS + RoPE stay on device at decode shapes and cur_pos advances on-device, with no per-step host readback. Default since #1604 (2.1x over the host-hybrid path: 10.94-11.06 vs 5.34 tok/s warm, Qwen3-0.6B batch 1). Exact 0 restores the pre-flip host-hybrid path for same-binary A/B; any other spelling keeps the default |
VT_TT_DECODE_CAPTURE |
unset (capture ARMED on the gated families) | Tenstorrent decode-graph capture on top of host-free decode. Default since #1625's flip landed in #2910, and exact 0 opts out and keeps gating the eager host-free arm. The default is EVIDENCE-SCOPED rather than global: it applies to the families with a committed captured-arm gate pair, while every other decode driver keeps the explicit opt-in until its own pair exists -- which is what Platform::static_graph_requires_opt_in() carves. This row said "capture declined" and "NOT the default" until 2026-09-05, on the strength of a deterministic multi-request hang; that hang was root-caused to #2669's short-chunk device KV push clobber, where prefill chunks below kPagedFillMinTokens ran the batched device update per token over one shared physical block and the last writer won, leaving captured decode attending a dead request's rows. Repaired, the captured multi-request battery is 16/16 and hang-free, and the captured arm is the fastest measured arm at 28.61 tok/s warm median against 12.21 for the eager host-free arm (same-binary A/B) |
VT_TT_DECODE_CAPTURE |
unset (capture ARMED on the gated families) | Tenstorrent decode-graph capture on top of host-free decode. Default since #1625's flip landed in #2910, and exact 0 opts out and keeps gating the eager host-free arm. The default is EVIDENCE-SCOPED rather than global: it applies to the families with a committed captured-arm gate pair, while every other decode driver keeps the explicit opt-in until its own pair exists -- which is what Platform::static_graph_requires_opt_in() carves. This row said "capture declined" and "NOT the default" until 2026-09-05, on the strength of a deterministic multi-request hang; that hang was root-caused to #2669's short-chunk device KV push clobber, where prefill chunks below kPagedFillMinTokens ran the batched device update per token over one shared physical block and the last writer won, leaving captured decode attending a dead request's rows. Repaired, the captured multi-request battery is 16/16 and hang-free, and the captured arm is the fastest measured arm at 28.61 tok/s warm median against 12.21 for the eager host-free arm (same-binary A/B) |
VT_TT_RELEASE_WARM_ROWS |
unset (keep all slots) | Tenstorrent-only decode-memory lever: after the engine's cold pre-warm, releases every slot whose consumer shadow has exactly this many rows — the warm pass's full-batch-shape activations, which the captured decode never reads. At 27B this frees ~14 GB of slot residency that otherwise filled the banks to 99 percent during generation. Row count must match the warm batch (e.g. 13104); unset keeps every slot |
VT_TT_WEIGHT_RESIDENCY |
off |
Tenstorrent-only weight-residency lever (supersedes the boolean VT_TT_BFP8_WEIGHTS): selects which device-side residency the staging seam converts an admitted weight to, parsed in include/vllm/config/tt_weight_residency.h. off (also unset, empty, or 0) is byte-identical bf16 staging; bfp8 uploads bf16 and typecasts to BFLOAT8_B at first staging, shrinking resident weight bytes to roughly half; bfp4 is RESERVED — it parses but every consumer refuses it by name ("not implemented") so a reserved value stays visible debt, never a silent fallback. An unknown value aborts loudly naming the string. Read fresh per call, so a test can flip it mid-process. See .agents/specs/tenstorrent-bfp-weight-residency.md |
VT_BENCH_PRETOKENIZE |
1 (on) |
Makes vllm-bench encode every prompt before its benchmark clock and admit token IDs, matching the pinned vLLM comparison frontend. Exact 0 restores timed string admission for same-binary A/B; unset, 1, and invalid spellings keep the safe default-on behavior |
VT_VK_FENCE_TIMEOUT_MS |
0 (block forever) |
Vulkan-only watchdog: caps how long a dispatch wait blocks in vkWaitForFences. The default is vLLM-identical stock behavior — a dispatch that never completes blocks forever at ~100% host CPU while a wedged GPU burns a core. Set a positive number of milliseconds to make the process ABORT instead, with a loud message naming the last dispatch submitted (BACKEND-VULKAN-FENCE-TIMEOUT, src/vt/vulkan/vulkan_context.cpp). 0, unset, and unparseable values all restore block-forever |
VT_VULKAN_DEVICE |
first suitable device | Forces the Vulkan physical device index. Required on a multi-GPU host to pin the intended device |
VT_KV_CACHE_F32 |
off (native KV dtype) | Forces the KV cache to fp32. A precision/diagnostic lever, at the cost of double the KV memory |
VLLM_CPP_VOCODER_DEVICE |
cpu |
Which device the shared 1-D BigVGAN vocoder core (vllm::vocoder1d — MiniMax-Music3, MiniMax-H3's audio VAE, LTX-2.5's audio VAE, IndexTTS-2.5) runs its convolutions on. It takes any device name vt knows (vt::DeviceTypeName — cpu, cuda, metal, vulkan, xpu, rocm, tenstorrent) and resolves it through vt::DeviceTypeFromName, so a provider registered for a new backend becomes reachable here with no edit. A name vt does not know, or one whose device has no registered vt::Conv1d / vt::ConvTranspose1d provider in this build, is REFUSED by name — never silently downgraded to the host, because a silent fallback means an operator who asked for a device never learns they did not get one. The transposed convolution is 88.5 % of MiniMax-Music3's acoustic-half profile and on scalar host loops a 45 s clip is a multi-hour decode, so this is the knob that decides whether that stage runs on the GPU. The two providers are BYTE-IDENTICAL — one f64 accumulator per output element walked in the same order, the host pinned -ffp-contract=off and the device kernel pinned with __dmul_rn/__dadd_rn — and tests/vt/test_ops_conv1d_general.cpp gates that with memcmp, not a tolerance. It still defaults to cpu: flipping four shipped audio models onto a device arm is not a default the row that ADDED the arm is entitled to set, and the flip is owed to the wiring row named in .agents/specs/minimax-music3.md §11.4. Requires a CUDA build — asking for cuda without one throws rather than falling back silently (#672) |
VLLM_CPP_MUSIC3_PROFILE |
unset (off) | Prints a per-STAGE wall-clock and RSS split for a MiniMax-Music3 request to stderr when it finishes. examples/minimax_music3_gen/main.cpp times only the checkpoint load and the whole of vllm_synthesize, and at a real duration those two numbers cannot say which of the six stages spent the time — four of the six are still host reference loops (.agents/specs/minimax-music3.md §11.1, §14.2), so the split is the question. Rows are LEAVES (they partition the request and are summed) or SPANS (they enclose leaves and are printed but never added), plus pure counters; unattributed is the glue between the leaves and is printed rather than spread silently over the measured stages. calls on denoise.dit_device counts denoise STEPS, so the DiT forward count is twice it. Accepts 1, true, on, yes (case-insensitive) — any other value, including a near miss like y, leaves it off, because an operator who mistypes should get a run with no table rather than a run whose meaning quietly changed. With it off no clock is read and no /proc file is opened. It is an attribution instrument and takes no GPU clock window (.agents/benchmarking.md), so its rows are a within-run split and must not be quoted as per-kernel or cross-box figures (#672) |
VLLM_CPP_MUSIC3_DIT_SPANS |
unset (off) | Splits VLLM_CPP_MUSIC3_PROFILE's single denoise.dit_device row into sixteen SPANS inside the flow-matching transformer forward — the host packing, the input projections, the timestep embedding, the rotary tables, nine per-layer spans and the output projections and readback — plus dit.seq_sum and dit.length_sum, the summed window geometry the split has to be read against. Ignored unless VLLM_CPP_MUSIC3_PROFILE is also on: it refines a table, it cannot create one. It is a SECOND flag rather than more detail on the first, and that is a property of the measurement rather than caution. The forward's operations are asynchronous on one stream, so a bracket that does not drain the queue times the LAUNCH — it would report every matrix multiply as free and charge its cost to whichever bracket happened to contain the next synchronize — so the drain is mandatory for the split to mean anything, and it perturbs the very total it splits. With this unset the forward is exactly the path every recorded denoise.dit_device figure was taken on (.agents/specs/minimax-music3.md §15.7, §20), so read the SPLIT from a run with the flag and the TOTAL from a run without it. The spans are printed and never summed, so sum(leaf), unattributed and every recorded stage total are unchanged either way. Same value grammar as VLLM_CPP_MUSIC3_PROFILE: 1, true, on, yes, case-insensitive, and anything else including a near miss leaves it off (#1542) |
VT_ENABLE_JUMP_FORWARD |
off | Opt-in to jump-forward constrained decoding (SGLang parity SW3): when a grammar/structured-output request reaches a state with exactly one valid next token, that token is emitted without a model step. Currently drives only the standalone driver (DrainForcedTokens); output-identical by construction (it fires only where the constrained sampler already has a single valid token), so it changes speed, never tokens. Off by default until the production scheduler splice (jumped-token KV recompute) lands. Set 1/true/on to enable |
Behavior of the GGUF weight loader (CPU path). See the quantization support table for the current format and backend coverage.
| Variable | Default | What it does |
|---|---|---|
VT_LAGUNA_FUSED_GATEUP |
off | Laguna only. Route the grouped routed-expert gate+up pair through the fused vt::MoeGateUpSwiGLUGrouped, which quantizes the activation ONCE where the default two-call arm quantizes it twice (QuantizeQ8KKernel was 12.4% of decode GPU time). OFF because the arm is not bit-identical: the shared epilogue computes g * (1/d) * u while Laguna's GateUpSilu computes (g/d) * u, a reciprocal-then-multiply against a divide, and a deterministic sweep puts 20.3% of values apart by at most 2 ULP (relative < 1e-6, sign never flips; pinned by test_laguna_fused_gate_up). Byte-identity would mean changing the shared op that DeepSeek-V4 is also gated against. The token gate ran on the real checkpoint (GB10, warm, order-balanced, both arms repeated) and the difference DOES move a token — on 6 of 6 prompts swept at 256 tokens, first diverging anywhere between token 2 and token 115: the streams diverge at position 2 and cascade, reproducibly, with both arms self-consistent. The lever is worth about 4% warm (7.85 to 8.19 tok/s, n=2 per arm, a direction rather than a ratified number), and the row refuses that trade — correctness over 4%. It stays off. Inert unless both expert towers share a block-quant dtype — a UD quant varies type per tensor, and a mismatch silently falls back to the two-call arm rather than refusing the load (#2061) |
VT_GGUF_KEEP_QUANT |
on when compute-in-quant is available | Keep GGUF weights compressed from file to matmul, byte-identical to the reference path. 0 disables it and expands to BF16 — which needs roughly 4x the file's compressed size in resident memory, not the file's size, since a Q4_K_M tensor's on-disk bytes are close to a quarter of its expanded BF16 size. This is a real precondition on a weight-staging device (CUDA/ROCm/...; CPU loads borrow the file's own pages either way): with VT_GGUF_KEEP_QUANT=0 and no other residency-shrinking flag active for the device (no VT_GGUF_KEEP_F16, no VT_GGUF_NVFP4_FP4), CheckDeviceWeightFit (gguf_device_fit.h, issue #1123) knows every tensor's residency exactly — RouteGgufTensor's decision is total once every keep flag is off — and REFUSES the load by name, citing the required and available bytes, before any allocation is attempted, rather than letting the checkpoint reach a raw allocator crash (hipMalloc/cudaMalloc: out of memory) on the first forward. See issue #1870 (reproduced on a 16 GiB ROCm card) and .agents/specs/gguf-device-fit-expand-policy.md |
VT_GGUF_NVFP4_FP4 |
on where the device can run the NVFP4 GEMM (CUDA; a CPU build expands) | The NVFP4 analog of VT_GGUF_KEEP_QUANT: keep an NVFP4 GGUF's weights in native fp4 residency and run kMatmulNvfp4, instead of expanding to BF16. 0 is the same-binary opt-out (expand to BF16); forced off under VT_CPU_REF so the oracle load stays byte-identical. See .agents/specs/gguf-nvfp4-native-compute.md |
VT_GGUF_NVFP4_W4A4 |
on (only meaningful when VT_GGUF_NVFP4_FP4 is on) |
Selects which of vLLM's two NVFP4 modes the fp4-resident weights compute in: on = true W4A4 (fp4 activations, using the GGUF's <stem>.input_scale sidecars, mirroring the sibling compressed-tensors container); 0 = W4A16 (BF16 activations over the fp4 weights). No effect when the fp4 residency is off |
VT_GGUF_KEEP_F16 |
on (when weights expand) | Keep F16 GGUF weights resident as F16 instead of promoting them to BF16, so the elementwise f16 GEMM computes on them directly (mirrors llama.cpp ggml_vec_dot_f16). On by default as a deliberate trade, and here is what it costs and buys. Measured on Qwen3.5-2B-UD-Q8_K_XL, aarch64, 20 threads, same binary: peak RSS 3.885 to 2.832 GiB, so 1.05 GiB less resident, bought with about 9% of prefill (224 to 204 tok/s) and about 1.4% of decode (TPOT 40.4 to 40.95 ms). Output tokens are byte-identical either way, so this is a memory-for-speed trade and never a quality one. The recorded reason for keeping it on is that a gigabyte decides whether a model fits at all, which is a harder failure than a slower prefill. Set 0 to opt out on the same binary: prefill and decode return to the faster arm and peak RSS returns to 3.885 GiB. That is the setting for a host with memory to spare. Decision and full A/B in .agents/specs/gguf-keep-quant-loader.md |
VT_GGUF_MMAP |
on when weights stay quantized | Keep the GGUF file mmap-resident instead of copying weight bytes into owned buffers, trading RSS for page-cache residency. This is what makes a checkpoint larger than host RAM loadable at all: the weights are borrowed from the mapping and never copied. Also settable as --offload-config's vllm_cpp.mmap.enabled; this variable overrides it. Forced off under VT_CPU_REF regardless, so the oracle load stays byte-identical |
VT_GGUF_PREFAULT |
on with mmap residency | Pre-fault the mmap-resident weight pages at load, trading a slower load for steadier first-token latency — without it those first-touch faults land in the TIMED prefill. Unset reads as ENABLED; 0/false/off disables it. Set it off for a model larger than memory, where prefaulting would read the whole borrowed tower to populate a page cache that cannot hold it. Also settable as --offload-config's vllm_cpp.mmap.prefault; this variable overrides it. This row said "off" until 2026-08-17 and the code has always defaulted it on (#1109) |
VT_H3_DROP_PAGES |
off (MiniMax-H3 device staging) | Set (to any value) to opt the mapping into page release while streaming the H3 DiT weights to the device, so the read-once file pages do not accumulate against the same unified pool the weights live in. Without it the DropSpanResidency calls are no-ops. Off by default deliberately, not by oversight: it was enabled once and the load was SIGKILLed early at only ~21 GB peak, which is not a memory ceiling, so it stays gated until that is understood rather than left on by faith |
VT_H3_NVFP4_LOWNIBBLE |
off = swap ON (MiniMax-H3 NVFP4 loaders) | The community lilcheaty/MiniMax-H3-NVFP4 checkpoint (metadata converted_by: "Star Ultimate Model Converter Pro") packs the two fp4 elements per byte HIGH-first (element 2i in the high nibble) — the opposite of the modelopt standard DequantNvfp4ToBf16 and the Marlin W4A16 path assume. By DEFAULT the three H3 NVFP4 loaders swap the two nibbles of every packed byte at load, turning the file's high-first bytes into the standard low-first (byte-verified against an independent oracle and the coherent FL2VA GGUF: reading low-first gives elementwise corr 0.000, high-first sign-agreement 1.000 over 115M+ weights). =1 disables the swap (the pre-fix low-first read) for A/B. H3-scoped: the shared modelopt NVFP4 arms (Laguna / DeepSeek-V4 / Qwen3-32B) are low-first and untouched |
Opt-in host-side residency cache for routed MoE experts on the CPU keep-quant
path, used when a mixture-of-experts GGUF is served from a borrowed (mmap'd)
weight tower larger than RAM — the Qwen3.5-family MoE decode path in
src/vllm/model_executor/models/qwen3_5.cpp. Borrowing is what makes such a
model fit; this cache is an attempt to make it fast, by serving each expert
slice out of a fixed array of equal-sized slots filled by one contiguous copy
instead of letting the kernel demand-page the mapping in router order.
Off by default, and honestly still experimental. The W4 measurement on
Qwen3.8-2.4T-A95B UD-Q1_0 (.agents/specs/expert-streaming.md, "W4 measured")
is token-identical output and a 4.5x TTFT improvement (3318 s to 733 s) with
no steady-state decode gain at all, because the slot is filled by copying
from the mapping and so still takes the page fault it was meant to avoid. Turn
it on to reproduce or to work on that; it is not yet a deployment
recommendation.
The two budget knobs are read once, when the cache is first constructed, and
only if streaming is on. Host memory reserved is slots x slot_bytes,
allocated up front and never grown — the engine prints the resolved values as
[expert-stream] ON slots=... slot_bytes=... resident=... GiB on stderr.
| Variable | Default | What it does |
|---|---|---|
VT_MOE_EXPERT_STREAM |
off | =1 serves routed expert slices from the bounded host slot cache instead of reading them straight out of the mmap'd tower. Read once per process, and only the FIRST character is examined: a value starting with 0, and an empty value, are off; anything else is on. Only the CPU path streams — on a device platform the expert slice is already device-resident and is served unchanged. Turning it on also disables the default-on grouped-MoE path (VT_QWEN35_GROUPED_MOE), which stages the whole tower and therefore cannot stream; the engine says so once on stderr rather than silently doing no streaming. Set VT_MOE_EXPERT_STREAM=0 to keep grouping. Also settable as --offload-config's vllm_cpp.expert_stream.enabled; this variable overrides it |
VT_MOE_EXPERT_STREAM_SLOTS |
64 |
How many expert slices stay resident. Parsed as a decimal integer; unset, empty, zero, negative and unparseable values all fall THROUGH to --offload-config's vllm_cpp.expert_stream.slots if that is set, and to 64 otherwise. So this variable overrides the config only when its value parses to a positive integer, and the startup override line says so on the same terms. Every slot acquired during a step is protected from eviction until the step ends, so a budget smaller than one step's working set exhausts the cache: those slices fall back to reading the tower directly, which is correct but slow, and is counted. Sized against a real model this wants to be large — the measured run used 8000. The config key vllm_cpp.expert_stream.slots, unlike this variable, REFUSES a zero or negative value at startup instead of silently keeping 64 |
VT_MOE_EXPERT_STREAM_SLOT_BYTES |
the LARGEST of the gate/up/down slices of the first MoE layer reached | Bytes reserved per slot, fixed for the process's life. Parsed as a decimal integer; unset, empty, zero, negative and unparseable values all fall THROUGH to --offload-config's vllm_cpp.expert_stream.slot_bytes if that is set, and to the default otherwise — this variable overrides the config only when its value parses to a positive integer. The default is the largest of the three slices rather than the first one taken, because a dynamic (UD) quant keeps down_proj at a higher precision than the gate/up pair and sizing from a gate slice then refuses the first down slice mid-decode. A slice that still does not fit is refused BY NAME (vt: expert stream: a slice of N bytes exceeds the slot budget of M; raise VT_MOE_EXPERT_STREAM_SLOT_BYTES) rather than truncated or silently routed back to the mmap path, so a streaming benchmark cannot quietly measure the mmap path instead. The config key vllm_cpp.expert_stream.slot_bytes, unlike this variable, refuses a non-positive value at startup |
VT_MOE_EXPERT_STREAM_STATS_EVERY |
16 |
How many decode steps between the PERIODIC expert-stream statistics line on stderr; 0 silences the periodic line only. Parsed as a decimal integer; unset, empty, negative and unparseable values all keep 16. The line is [expert-stream] steps=N hits=H misses=M evictions=E fills=F bytes=B exhausted=X advised=A. Exactly one FINAL line is printed when the process ends, whatever this is set to and whatever the run did, including steps=0, for as long as the lane built a store. That is the line to read, and it exists because the row's first published decode figure was measured on a cache that had switched itself off partway through the third token while nothing in the run could say so. steps == 0 or exhausted > 0 means the lane is not streaming, whatever the startup line claimed. Absence of the final line means either that no store was ever built — in which case the [expert-stream] ON ... banner is absent too, and the lane was never reached — or that the process did not run its static destructors (a crash, a signal, _exit). A fourth shape exists but no shipped command can produce it: the line is printed once per process, and an internal test seam that flushes it mid-run takes that one print. docs/USAGE.md tabulates all four. Environment-ONLY by decision: it changes a diagnostic cadence rather than what the process reserves, so it is the instrument and not the configuration, and --offload-config refuses it as an unknown key rather than accepting and dropping it |
VT_DEVICE_WEIGHT_BUDGET_BYTES |
the device's own probe (cudaMemGetInfo total on CUDA; UNKNOWN, i.e. no check, everywhere else) |
Also settable as --offload-config's vllm_cpp.device_fit.weight_budget_bytes; this variable overrides it, and a value outside this variable's digits-only grammar falls THROUGH to that key before it falls through to the probe. Overrides the device memory pool that a GGUF's staged weight bytes are compared against at LOAD time (issue #1123). A GGUF whose weights cannot fit is refused by name during the load instead of dying on the first forward with vt cuda: cudaMalloc: out of memory — Qwen3.8-2.4T-A95B UD-Q1_0 (369.96 GiB) reached a serving state on a 119.631 GiB GB10 after 26 minutes and then died mid-stream, because the larger-than-memory lane that makes it fit is HOST-ONLY. Set this LOWER when something else lives in the pool, or HIGHER (or 0) to suppress the refusal and get the late failure back — it does not make the model fit. Parsed as decimal digits only: a value with a sign, a space or trailing garbage is IGNORED and the next tier decides — the config key if it is set, otherwise the probe — because reading a typo as 0 would silently disable the guard. Compared against the pool TOTAL, not the free bytes, so the verdict does not move with contention. The bound counts WEIGHTS only, never the KV cache, activations or the driver context, so a checkpoint just under the pool still passes and can still fail later. It can also count a little too MUCH: a tensor present in the file that this load will not stage — the MTP / nextn block on a load with no speculator, 8.33 GiB of the measured 369.96 GiB checkpoint — is still in the sum, so a budget in that narrow window refuses a weight set that would have fitted; raise this value if you land in it (#1136). Inert on every platform that does not stage weights, which today means everything except CUDA — including every --device cpu load, and including ROCm, Vulkan and Metal, which read the mapping where it lies and have no staging allocation to fail |
VT_PLACEMENT_DUMP_MOE_LAYER |
0 |
Which decoder layer VT_PLACEMENT_DUMP_MOE writes. Selectable because a fixed layer 0 makes a --fit comparison VACUOUS: the resolver places TRAILING layers, so a gate dumping layer 0 compares an UNPLACED layer against an unplaced one and reports a perfect NMSE of 0 over bitwise-identical values — a pass that measured nothing, because the dump region and the placement region did not intersect. Point it at a layer the placement actually covers (for --fit, one of the last N; for cpu_moe, any). Latched once per process alongside the path, and parsed with atoll: unset, empty, negative or unparseable all mean layer 0, so every existing invocation is byte-identical. A layer index the model does not have simply never matches, and the gate then sees a MISSING dump rather than a silently empty comparison |
VT_PLACEMENT_DUMP_MOE |
unset | Names a file the MoE placement seam writes LAYER 0's block output to, ONCE, on whichever path ran — placed or unplaced. It exists so the placement correctness gate can compare a placed run against an unplaced one NUMERICALLY. A token comparison cannot serve here and is not merely noisy: a placed layer computes its experts with the CPU MoE kernels rather than the accelerator's, and this project's cross-device bar for reducing arithmetic is NMSE <= 5e-4 rather than bitwise equality (tests/vt/test_backend_cross_device.cpp, where CPU is the oracle), so greedy decode amplifies any perturbation inside that bar into a different token — measured on GB10, both arms answer the prompt correctly and diverge mid-sentence, exactly as llama.cpp's -ncmoe does. Values are written as one %.9g float per line, widened from bf16, which is exact. FIRST matching call only: a decode writes the layer once per step and the gate compares one step against one step, so appending every step would compare arms that have already diverged in tokens and no longer share an input. INERT unless set — one latched getenv, no allocation, and no copy on the placed path because the staging buffer already holds host bytes. A path that cannot be opened is reported on stderr rather than skipped silently, because a gate that finds no file would otherwise compare nothing and read it as agreement |
VT_CPU_MOE |
off | Places routed experts on six architecture families — see the note below the table. Our -cmoe. Ask for every routed expert to run on the CPU while attention and the dense projections stay on the device. Sugar for ONE override carrying llama.cpp's own LLM_FFN_EXPS_REGEX, and the install line reports the count it expanded to. Also settable as --offload-config's vllm_cpp.placement.cpu_moe; this variable overrides it in both directions, so VT_CPU_MOE=0 turns off a configured true |
VT_N_CPU_MOE |
unset | Places routed experts on six architecture families. Our -ncmoe N. Run the routed experts of the FIRST N layers on the CPU. Sugar for N overrides, one per layer index in [0, N). 0 is legal and places nothing, matching the flag this mirrors, so unlike the slot counts it is read with a non-negative grammar and a 0 DOES override a configured value. A negative or unparseable value is ignored and the config decides. Also vllm_cpp.placement.n_cpu_moe |
VT_PLACEMENT_OVERRIDES |
unset | Places routed experts on six architecture families. Our -ot, and the general form the two above are sugar for. Takes llama.cpp's own grammar, <pattern>=<device>,..., so a command line ported from there keeps working; the split is on the FIRST = because a regex may contain one and a device name may not. <device> is a vt::Backend name (cpu, cuda), NOT a ggml buffer-type spelling. ORDER MATTERS and is never sorted: the list is scanned first-match-wins, so a narrow entry before a broad one wins. A malformed value is ignored and the config decides. Also vllm_cpp.placement.overrides, which takes an array of {pattern, device} objects and refuses a bad regex at startup |
VT_PLACEMENT_FIT |
on | Places routed experts on six architecture families. Our --fit, and it is ON BY DEFAULT because llama.cpp's is (fit_params = true, common/common.h:468 @ b10451): upstream's headline behaviour is that a model too large for the device just runs, without the operator configuring anything, and shipping the mechanism switched off would give the machinery without the experience. Asks for the placement to be resolved from the device budget instead of stated by hand, and resolves to a whole number of TRAILING layers whose routed experts move to the CPU. Three things bound the default. It fits against the platform's STATIC TOTAL, not live free — deliberately, because include/vt/backend.h records that a load-time budget must not be a function of contention, so the same model with the same flags resolves the same way whatever else is running; the cost is that a box with memory already in use can be told a model fits when it does not. An UNKNOWN budget (0, which is what a platform that does not probe one reports) places NOTHING rather than everything. And it applies to GGUF only: a safetensors checkpoint's weight footprint is not known where the placement must be installed, so the resolver is INERT there when defaulted — one line on stderr and no placement — and REFUSES BY NAME only when you asked for it explicitly. Whole layers only, so unlike upstream it cannot spill half a layer; when it places, it prints the budget, the footprint, the layer count and the bytes. =0 turns it off. MUTUALLY EXCLUSIVE with a manual placement rather than merged with it: a config document setting both is refused at startup by name, mirroring llama.cpp's own refusal. Also vllm_cpp.placement.fit |
VT_DSV4_EXL3_HOST_BUDGET |
on (the refusal armed) | Guards the DeepSeek-V4 EXL3 loader arm (a checkpoint whose quantization_config.quant_method is exl3), which COPIES the TP1-coalesced trellis tower into HOST owner buffers AND dequantizes the carried-* non-expert tensors into a host-float tower beside it — the real artifact's trellis alone is ~83.5 GiB and its carried half adds ~29 GB of f32. The loader prices the trellis tower from the layers it has already built, adds the carried tower it has MEASURED, and refuses the load BY NAME, before the allocation, when the sum exceeds the budget; on a unified-memory box an over-commit reboots the machine instead of failing a malloc, so a late failure is not an option. The budget is /proc/meminfo's MemAvailable, and an UNKNOWN budget (0, which is what an unreadable /proc/meminfo yields) never refuses. Set this to 0 to disarm the refusal and proceed in the SAME binary — no rebuild. Read once per process, and its parse is NARROWER than the flag rule at the top of this page: the refusal is off only when the value's FIRST character is 0, so unset, empty and every other value leave it ON and false / off do NOT disable it. Two caveats make this override necessary rather than convenient. MemAvailable is an ESTIMATE that ignores swap and under-counts some reclaimables, so the guard can refuse a tower this host would in fact have held. And inside a container it reports the HOST's figure rather than the cgroup limit, so a memory-capped container gets NO protection from this guard while the logged budget names a pool the process cannot draw on. A device-resident / per-layer-streaming destination, which would remove the host copy this guards, is owed to MODEL-DSV4-EXL3 W2. |
The Kolibri-1 CPU decode path interleaves the weight-arena allocation and the
forward pass over every NUMA node the process may use
(MPOL_INTERLEAVE, src/vllm/model_executor/models/kolibri1_numa.h, wired at
kolibri1_weights.cpp in LoadKolibri1Weights). Multi-node hosts otherwise
keep the arenas on the loading node's memory, and 32-thread decode then reads
most weight bytes remotely. Measured on the 4-node aarch64 host
(docs/bench-evidence/kolibri1-numa-interleave-20261009.md): 4.19 to 4.97
decode tok/s at 32 threads with the 16 GiB dequant cache (+18.6%), prefill
−14%, golden decode chain unchanged. The guard saves and restores the caller's
mempolicy, and it is inert on single-node hosts and on Windows. The lever
composes with the dequant cache; with the cache off and 8 threads its effect
collapses to ~2%, because that path is allocation-bound.
| Variable | Default | What it does |
|---|---|---|
VT_KOLIBRI1_NUMA_INTERLEAVE |
on | =0 (exactly) disables the guard entirely: the arenas allocate under the process's default policy and nothing touches the mempolicy. Unset or any other value leaves it on. Read on every guard construction (uncached), so a test binary can flip it mid-process; the call rate is per load and per cache miss, never per element. This is the A/B off switch and the first thing to set to 0 when a NUMA-policy defect is suspected on a multi-node host |
Default-on fast paths, each with an off switch. They exist so a suspected kernel
bug can be bisected without a rebuild: set the variable to 0 to fall back to the
portable/reference path. In normal operation leave them unset.
| Variable | Default | Off switch falls back to |
|---|---|---|
VT_ASYNC_RUNNER |
on | Synchronous model runner (no async/overlap execution) |
VT_DSV4_EXL3_FUSED_MOE |
on | The per-expert vt::Exl3Gemm loop, for a DeepSeek-V4 checkpoint whose quantization_config.quant_method is exl3 (MODEL-DSV4-EXL3). ON, the layer's routed experts go through ONE vt::Exl3MoeMlp call that gathers each expert's tokens, runs gate/up, the activation and down, and scatter-adds the weighted result; OFF, each (token, expert, projection) is its own vt::Exl3Gemm — 3 * topk * T calls per layer against one. The two arms compute the same algebra by different roundings, because the fused arm keeps the intermediate in fp16 through the activation the way the upstream kernel does while the loop widens to f32 and back, so they agree only within the row's tier-4 bound (2.0e-2 relative RMS) and not bit for bit. That is what this switch is for: a decode that looks wrong can be bisected onto one arm or the other in the SAME binary. It is NOT a correctness escape hatch, and it does not turn the loop off: an expert with more tokens than the fused kernel's staging buffers hold (128) is declined by the fused arm and runs per-expert whatever this is set to, which is upstream's own arrangement. Parsed like the other VT_DSV4_EXL3_* knob rather than by the flag rule at the top of this page: the fused arm is off only when the value's FIRST character is 0, so unset, empty and every other value leave it ON and false / off do NOT disable it. Inert on any model that is not an EXL3 DeepSeek-V4 |
VT_ASYNC_SCHED |
on | Synchronous scheduling (no scheduler/execution overlap). The documented first-line workaround for a suspected scheduling bug |
VT_ASYNC_DEVICE_MIRROR |
on | The CUDA device-resident sampled-token mirror (ENG-ASYNC-SCHED W4): the async serving loop's sampled ids stay on the device instead of round-tripping the host. On a DISCRETE GPU this removes the host-fallback main-stream Synchronize. On the INTEGRATED GB10 it moves last_sampled_tokens off the host array update_states' condense reorders (a drain move, not a removal). DEFAULT ON since the 2026-08-06 correctness flip (ROW-SERVE-ASYNC-LLM): the OFF path's host-array combine writes step.input_token_ids on the main queue while the decode graph reads it on the CPU without a sync, so async batch-1 greedy decode nondeterministically degenerates into token-0 garbage; the mirror routes the ids into the embed on-queue and fixes it. =0 is the rollback to the (racy) host-array path. Speed-neutral (c16 0.999x). No effect on CPU or the sync LLMEngine |
VT_ASYNC_EXECUTOR |
off (opt-in) | =1 enables Option A: the decode-graph per-step input H2D staged OUT of the captured replay (the c16/c32 overlap unlock, ENG-ASYNC-SCHED). The Qwen3.5 MoE/dense decode-graph drivers give each padded-size slot PERSISTENT device input buffers the captured graph reads and PINNED host staging; per step the input H2D is enqueued on the main queue BEFORE ReplayGraph and an input-staged event is recorded right after it, so the next same-slot Refresh waits only that tiny copy, never the GPU tail (the faithful vLLM _prepare_input_ids/synchronize_input_prep structure, states.py:64). The 2-slot parity ring is retained (the depth-2 loop enqueues sample(i-1) after forward(i), so persistent logits must double-buffer). The runner skips the depth-2 pre-forward Synchronize whenever the previous step's logits are a non-owning graph-slot view. Default OFF routes through the single-slot baked-H2D driver with the drain intact — byte-identical to production. No effect on CPU or the sync LLMEngine. =1 roughly doubles the captured decode-graph memory (logits-dominated) and adds small pinned host + persistent device input buffers per slot |
VT_GDN_CHUNKED |
on | The exact SEQUENTIAL gated-delta-rule recurrence, instead of vLLM's chunked WY decomposition, for GDN prefill on every device that has both arms (CPU and CUDA today; Tenstorrent has only the chunked one). Decode is sequential either way and is unaffected. ON is the mirror: vLLM runs the chunked decomposition for all GDN prefill and has no sequential branch in the layer at all. =0 selects an arm that is measurably NEARER the exact recurrence (1.15e-08 from it, against vLLM's own 2.29e-04) and correspondingly FURTHER from what vLLM computes, so it changes the token ids a model emits -- it is a bisect switch and an accuracy reference, not a quality setting. The dtype decides first: on f32 inputs the sequential arm runs whatever this is set to, because vLLM refuses f32 chunked on both of its implementations (chunk.py:213-215 asserts, csrc/cpu/sgl-kernels/fla.cpp:2205-2207 type-checks bf16), so at f32 the sequential recurrence IS the mirror. Parsed like VT_DSV4_EXL3_FUSED_MOE rather than by the flag rule at the top of this page: the chunked arm is off only when the value's FIRST character is 0, so unset, empty and every other value leave it ON and false / off do NOT disable it. That spelling is deliberately unchanged -- it is the off-value recorded in four evidence files (KERNEL-GDN-CHUNKED-MIRROR, #2612) |
VLLM_CPP_CUDAGRAPH |
on (CUDA) | Eager launches instead of a captured CUDA graph |
VLLM_CPP_ROCM_STATIC_GRAPH |
off (ROCm) | T2b per-arch static-graph opt-in: on gfx1100 the captured decode-graph path engages by default (evidence arch); every other ROCm arch stays eager until it carries its own captured-arm evidence. Setting =1 opts IN on those arches deliberately; =0 forces eager even on gfx1100. Honors VLLM_CPP_CUDAGRAPH=0 upstream of the arch check |
VT_CUDA_GRAPH_DEDUP |
off (opt-in) | =1 folds captured graphs that share a node topology onto ONE graph executable, re-pointing it with cudaGraphExecUpdate / hipGraphExecUpdate instead of instantiating a second (ENG-CUDAGRAPH-DEDUP, #1162). Today a model holds one executable per padded decode bucket — 7 at max_num_seqs=32, 11 at 64 — times eight capture drivers. A MEMORY and capture-time change, not a throughput one: a deduped replay launches the same nodes, and each candidate fold is probed with the real driver update on a throwaway executable before it is honoured, so a capture the driver refuses simply keeps its own executable. It logs vt graph dedup: captured N graphs, deduped to G execs per capture so the ratio is readable. Default OFF and it stays off until measured: a workload that alternates padded buckets every step pays one cudaGraphExecUpdate per switch, and the device byte-identity A/B that would price that is still owed. Unset leaves the capture path byte-identical to the pre-dedup one |
VT_CUDA_GRAPH_DEDUP_COARSE_KEY |
off (opt-in) | Only read when VT_CUDA_GRAPH_DEDUP=1. =1 drops the launch dimensions, the memcpy extent and the memset width from the dedup signature, keeping the kernel function pointers, sharedMemBytes, the memcpy kind, the memset element size and height, and the whole topology (ENG-CUDAGRAPH-DEDUP, #1226). Why it exists: the device A/B of 2026-08-18 measured the fold NEVER happening -- N == M in every VT_CUDA_GRAPH_DEDUP=1 cell -- because the padded batch dimension sits in exactly those dropped fields, so two decode buckets never shared a key and cudaGraphExecUpdate was never attempted. Why it is safe to try: the signature is only a lookup key. Every candidate fold is probed with the real driver update on a throwaway executable first, and a refusal gives that capture its own executable, so a coarser key costs a wasted probe and never a wrong replay. The registry logs vt graph dedup: captured N graphs, deduped to M execs (probes=P refused=R) per capture and prints the driver's own reason on each refusal, which is what separates "the key never grouped" from "the driver said no". The process announces vt graph dedup: key mode = COARSE or = EXACT once. PENDING_VERDICT |
VLLM_CPP_DENSE_DECODE_GRAPH |
on (CUDA dense) | Non-graphed dense decode |
VLLM_CPP_QWEN3_DENSE_DECODE_GRAPH |
on | Routes pure-decode steps for the SHARED pure-dense forward (Qwen3DenseModel, i.e. Qwen3 / Llama / InternLM3 / Mistral / InternLM2 ForCausalLM) through the captured decode CUDA graph. DEFAULT-ON since QUANT-CT-MXFP4-MARLIN-STRUCT (parity-enabler; banks ~+1.3% TPOT@c8 by collapsing the eager inter-kernel launch gap). Token-exact with eager — =0 opts out to the byte-identical eager decode; gated dgx SACRED test_qwen3_paged_engine 184/184 (graph ON==OFF, Qwen3-0.6B near-tie + 4B), async 82/82, Qwen3-8B-MXFP4 #44 smoke 3/3, Qwen3-32B-NVFP4A16 142/142. Honors VLLM_CPP_CUDAGRAPH=0 |
VT_MOE_FUSED_W13 |
on | Runs a dense/shared MLP gate+up PAIR as ONE Marlin GEMM over the N-concatenated [2I,H] operand + SiluAndMul (vLLM's merged gate_up_proj structure), instead of two split GEMMs + MoeSiluMul. Generalized NVFP4→MXFP4 in QUANT-CT-MXFP4-MARLIN-STRUCT (drops the classic-dense Qwen3-8B-MXFP4 decode from 180→144 marlin GEMM/step = vLLM-structural parity). =0 is the split A/B fallback. Numerically equivalent to split (the fused fp32 split-K reduce regroups by 1 bf16 ULP; token-exact vs the oracle — #44 fused==split 3/3, unit test_linear_method 99.9% bit-exact) |
VT_MARLIN_E1_PAR1 |
off (opt-in) | =1 clamps the single-expert (num_experts==1) DECODE marlin grid to sms×1 (48 CTAs on GB10) instead of the persistent sms×par grid determine_exec_config auto-picks (par=3 → 144 CTAs), matching vLLM's dense marlin tile-per-CTA count. Measured same-tool in-model (Qwen3-8B-MXFP4 c8): marlin 17,463→16,512 us/step (−5.4%, per-call 121.3→114.7 vs vLLM 113.1 = near-parity), TPOT 37.22→36.23 ms, token-exact on 8B-MXFP4. DEFAULT OFF: par regroups the fp32 C_tmp reduce, so the E=1 output differs by 1 bf16 ULP — on the 64-layer Qwen3-32B-NVFP4A16 that accumulates into a strict-token flip vs its committed SACRED anchor (test_qwen3_32b_nvfp4a16_paged_engine REQUIRE :344). Real MoE (num_experts>1) and prefill (thread_m_blocks>1) are untouched (byte-identical) |
VT_MARLIN_DENSE |
on (opt out =0) |
=1 routes the E=1 dense NVFP4/MXFP4 projections (dense_nvfp4_gemm.h MatmulNvfp4MarlinD/GateUpFusedMarlinD) through vLLM's OWN dense marlin GEMM (vt::MarlinDenseGemm) instead of the single-expert MoE-marlin route. The dense kernel is direct-A + tile-per-CTA with vLLM's own dense fp32-C_tmp reduce, so at M≤8 it runs the sms-wide (48-CTA) grid VT_MARLIN_E1_PAR1 targets but WITHOUT that flag's par-regroup ULP — the byte-preserving fix for the strict-32B token flip (row KERNEL-MARLIN-DENSE-PORT). Reuses the same marlin resident + workspace (shared marlin_permute repack). DEFAULT ON since efa6e40d (#57, "dense route beats MoE on every axis"); this page said OFF until 2026-08-10. CUDA-only (needs VT_MARLIN_NVFP4) |
VT_MARLIN_DENSE_PAIR |
on (opt out =0) |
Same dense route for the ONE fused shared-expert gate_up sink (SharedGateUpFusedMarlinD), which VT_MARLIN_DENSE did not cover and which therefore still ran the single-expert MoE-marlin route. Measured at c8 that was 20320 launches per leg (one per layer per step) of a <128,1,8,4,m_block_size_8=false> configuration the pinned vLLM never launches, 5.4% of GPU time. Dense route A/B (same binary, 3 reps/arm, order-alternated): +1.31% at c8, +1.38% at c4 on 35B-A3B NVFP4, non-overlapping bands against a +0.29% control. SACRED gates unmoved (35B 315/315, 27B 235/235). CUDA-only (needs VT_MARLIN_NVFP4) |
VT_DENSE_MARLIN_GATEUP |
on (opt out =0) |
Routes the dense MLP's W4A16 NVFP4 gate/up pair through the SAME fused Marlin gate_up GEMM the shared expert already takes (SharedGateUpFusedMarlinD: one [T,H]x[2I,H] operand + SiluAndMul), instead of two split Marlin GEMMs + MoeSiluMul. vLLM's dense Qwen3.6 MLP is one MergedColumnParallelLinear gate_up_proj, so ONE GEMM is the mirrored topology and the split pair is our divergence — measured 193 Marlin calls/step against the oracle's 129 on the identical batch-1 27B decode (#365; ~29% of a +4.40 ms/step gap, so it does NOT reach parity on its own). Selected only when the pair is NOT true-W4A4 (the CUTLASS merged path already covers that), NOT MXFP4 on either half, n/k/group_size/format match, scale2 is EQUAL on both shards (the merged operand carries one global scale), and VT_BF16_GEMM_OUT / VT_MOE_FUSED_W13 / VT_NVFP4_MARLIN are on. The MXFP4 refusal is NOT redundant with the format-match term (which only requires the two shards to AGREE): the fused entry point this lever reaches, SharedGateUpFusedMarlinD, is NVFP4-only — merged scale grid sized at K/16, NVFP4 combined-factor + global processing, group_size = 16/mxfp4 = false pinned in the GEMM args — so a group-32 E8M0 pair would be read as group-16 fp8-e4m3, verbatim the defect recorded RED-first at tests/vllm/model_executor/layers/test_linear_method.cpp:185-201. Latent, not live: no dense loader sets is_mxfp4 (LoadNvfp4AnyNaming), so refusing it moves no shipped configuration. Shipped OFF while row PERF-27B-DENSE-MARLIN-GATEUP held; FLIPPED ON once its same-binary A/B measured a win — interleaved 4 reps/arm on nvidia/Qwen3.6-27B-NVFP4@0893e160 (GB10), toggle the only variable, caches dropped between arms: +2.12% at c1 (12.0823 vs 11.8313 tok/s) and +1.70% at c8 (83.6186 vs 82.2217), with COMPLETE SEPARATION at both concurrencies (c1 min fused 12.0203 > max split 11.8729; c8 min fused 83.2511 > max split 82.4959) and a byte-identical 64-token greedy continuation on both arms. CUDA-only (needs VT_MARLIN_NVFP4) |
VT_SHARED_DOWN_BF16 |
on (opt out =0) |
The shared-expert down_proj emits bf16 instead of f32. The Marlin GEMM already produces bf16 and BOTH consumers (SharedExpertGate, MoeCombineGate) widen in-kernel and re-round through bf16 on store, so the f32 form wrote and re-read a whole [T,H] buffer for a value it already had. Bit-identical (widening bf16 to float is exact): SACRED 315/315 + 235/235 on BOTH arms with assertion counts unchanged, and the warm-server greedy probe is byte-identical across arms. Drops one CastF32Kernel launch per layer per step (CastF32 was 3.1% of the 35B decode step). Measured +2.05% at c8, +0.79% at c4 on 35B-A3B NVFP4, non-overlapping bands. CUDA-only |
VT_FP8_CUTLASS_SMALL_M |
on (opt out =0) |
The two SMALL-M rungs of the per-tensor FP8 CUTLASS GEMM's sm120 tile ladder (vt::MatmulFp8Cutlass, reached by default on VT_DENSE_CUBLASLT_FP8=0 and as the cuBLASLt lane's no-heuristic fallback). =0 restores the two-way ladder this tree shipped before #1866, where every M <= 256 took the 64x64x128 tile. Upstream's ladder is four-way (csrc/libtorch_stable/quantization/w8a8/cutlass/c3x/scaled_mm_sm120_fp8_dispatch.cuh:155-179 at the pin 5559679229): M<=16 → 16x64x128 with a 16x32 EpilogueTile, M<=32 → 32x64x128 with 32x32, M<=256 → 64x64x128, else 128x128x128. We ported the file and dropped the two small rungs as "perf-only for tiny M" — and tiny M is decode, so a batch-1 step computed a 64-row tile for one row and the 9-row spec-decode verify computed one for nine. Expected bit-identical, and not yet measured: all four configs use TileShape_K = 128 with no split-K or stream-K, so every output element accumulates over K in the same order in the same f32 accumulator whichever M-tile owns it; the M-tile choice moves which CTA computes a value, not the sequence of adds that produces it. The CUDA compile is measured green by cuda-fat-build (120a;121a, VLLM_CPP_CUTLASS_FETCH=ON) on head d9bf525c0, run 32802716762 -- the ONE job, not the run, whose own conclusion is cancelled; nothing since d9bf525c0 changes compiled CUDA and the current head's own cuda-fat-build is PENDING; the token gates on both arms and the VT_DENSE_CUBLASLT_FP8 decode A/B this makes fair are OWED (.agents/specs/perf-fp8-small-m-dispatch.md), so no speed number is claimed. CUDA-only, sm_12{0,1}a |
VT_MM_DECODE_EAGER |
off (graph on) | Set to 1 to force the eager per-step multimodal (Qwen3.6-27B image/video) decode instead of routing it through the captured dense decode graph. Rollback / A-B knob; the graphed path is token-exact with the eager path |
VT_KIMI_DEVICE_COMPUTE |
off (opt-in) | =1 routes the Kimi-Linear-48B-A3B runner path (KimiLinearModel::ForwardDevice) through the W7 DBuf-resident device COMPUTE (ForwardDeviceCompute, the whole KDA/NoPE-MLA + MoE hybrid over pooled DBufs via the shared vt:: ops) instead of the default W6 host-reference compose. Default OFF keeps the CPU-verified host-ref-compose seam as production until the device compute is GPU-verified against the SACRED oracle; the device compute is CPU-gated (test_kimi_linear_forward, device==W2 reference within f32-accumulation tolerance, greedy-token-identical) but its GPU numerics are a NAMED pending. The flag exists so the device path CAN be exercised as the runner path for that verification |
VT_KIMI_BF16_RESIDUAL |
off (opt-in) | =1 carries the Kimi-Linear device-compute residual stream in bf16 like vLLM's fused_add_rms_norm (residual/block-outputs bf16, RMSNorm variance over the f32 pre-store sum), via in-place f32→bf16→f32 rounds. W7-speed STRICT-lever A/B (spec §14). Default OFF → byte-identical. MEASURED: alone net-zero; WITH VT_KIMI_BF16_ISLANDS → 120/128 (best, still a near-tie, NOT STRICT) |
VT_KIMI_BF16_STREAM |
off (opt-in) | =1 carries the Kimi-Linear device-compute inter-layer residual stream in bf16 END-TO-END (STRUCTURAL bf16 DBufs for hidden/residual/normed-hidden/block-outputs via vt::FusedChain(kFusedAddRmsNormStd), mirroring deepseek_v2.cpp:479-615), SUPERSEDING the partial VT_KIMI_BF16_RESIDUAL in-place-rounding knob; elides the per-GEMM CastBf16 on residual-fed projections. MEASURED NEGATIVE / REFUTED (spec §20, GB10 full 48.9B 128-gate): REGRESSES 122→4/128 — the bf16 residual rounding DESTABILIZES the KDA recurrence into degenerate repeat loops (the §14/§15 pathology, confirmed structurally), with NO speed win (18.9→19.8, within noise). The p7-STRICT close is NOT reachable by this or any residual-precision lever; kept as a documented-measured-negative A/B (default OFF → byte-identical to the f32 stream), per the VT_KIMI_DEVICE_MLA/VT_KIMI_ISLAND_F32ACC precedent |
VT_KIMI_BF16_ISLANDS |
off (opt-in) | =1 rounds the Kimi-Linear host-fallback island INPUTS (KDA q/k/v/g1/beta, NoPE-MLA q/kv/kpe) to bf16 (RNE) before the recurrence/softmax, toward vLLM's GDN-Triton/FA2 kernel precision. W7-speed STRICT-lever A/B (spec §14). Default OFF → byte-identical. MEASURED best config paired with VT_KIMI_BF16_RESIDUAL (106→120/128) |
VT_KIMI_ISLAND_F32ACC |
off (opt-in) | =1 computes the Kimi-Linear island recurrence/softmax in f32 accumulation (not f64). W7-speed A/B knob, MEASURED NEGATIVE (91–106/128; kept as a documented-negative A/B, spec §14). Default OFF → byte-identical |
VT_KIMI_DEVICE_KDA |
off (opt-in) | =1 runs the Kimi-Linear KDA per-k-channel gated-delta RECURRENCE through the net-new device op vt::KdaGatedDeltaRule (cuda_gdn.cu KdaScanKernel, the per-channel-decay g[T,H,D] variant of GDN's per-head-scalar decay; FLA fused_recurrent_gated_delta_rule_fwd_kernel IS_KDA=True) instead of the f64 host recompute. The principled path to STRICT AND the speed lever (spec §14/§15): the recurrence runs vLLM's actual f32-on-bf16 arithmetic on device rather than a host f64 recompute that is more precise than vLLM and coin-flips near-ties. Requires VT_KIMI_DEVICE_COMPUTE=1; the elementwise decay gate + beta stay host. Unit-gated (test_ops_kda_recurrence: broadcast-g == vt::GdnPrefill bit-identical, distinct-per-channel vs f64 ref, CPU↔CUDA parity). Default OFF (parity-enabler: flip only with the token gate green) → byte-identical |
VT_KDA_CHUNK_TRITON |
on (default) | =0 restores the recurrent form inside the KDA chunk-prefill op, bypassing the Triton-AOT cubin path when the geometry is otherwise eligible (cuda_gdn.cu). An A/B escape hatch for isolating the AOT kernels from the recurrence; ON is the shipped behaviour |
VT_KIMI_PAGED_KDA_CHUNK |
on | =0 disables the CHUNKED KDA prefill (vt::KdaChunkPrefill, vLLM's prompt path kimi_gdn_linear_attn.py:141) inside the ROW 7 shared-paged-runner forward (KimiLinearModel::ForwardPaged) and falls back to the recurrence for fresh prefill requests — the fold-identity A/B. Decode and continuing (has_initial_state) prefills always use the recurrence, exactly as vLLM. Default ON is the §19 Gate-A-winning regime |
VT_KIMI_PAGED_MLA_FA2 |
on | Default ON (GB10-ruled 2026-08-07, kimi-linear.md §21): the ROW 7 paged-runner forward routes the 7 NoPE-MLA layers through mla::ForwardMlaAttentionBlock (vLLM's ACTUAL absorbed-MQA decode / FA2 prefill over the paged latent cache, identity-RoPE, scale qk^-0.5, load-time AbsorbKvBProjBf16) — measured on the real 48.9B it reproduces the golden's 122/128 near-tie profile EXACTLY (p0-p6 16/16, p7 10/16, got-string byte-equal to the CLI reference). =0 selects the DIAGNOSTIC exact f64-softmax-island arm (the CPU fold-identity vehicle; GB10 111/128 = the §19 GPU M-tiling near-tie regime, recorded, not a paging bug) |
VT_KIMI_DEVICE_MLA |
off (opt-in) | =1 routes the Kimi-Linear 7 NoPE-MLA layers' causal-softmax core through the shared device op vt::Attention (f32 online max-subtracted softmax, value zero-padded qk_nope+qk_rope=192 vs v=128) instead of the f64 host MlaSoftmaxIsland. Requires VT_KIMI_DEVICE_COMPUTE=1. W7-speed STRICT-lever A/B (spec §16, #107), MEASURED NEGATIVE — vt::Attention's f32 online-softmax is the right math but a DIFFERENT reduction order than vLLM's FA2, so it coin-flips near-ties (full 48.9B gate 122→109/128 AND 4.24→3.89 tok/s). Kept as a documented-MEASURED-NEGATIVE A/B knob; the STRICT path is the ACTUAL paged FA2 mla::ForwardMlaAttentionBlock, not this approximation. Default OFF → byte-identical |
VT_KIMI_DEVICE_KDA_CHUNK |
off (opt-in) | =1 processes the PROMPT-length Kimi-Linear KDA through the CHUNKED prefill kernel family (vt::KdaChunkPrefill: the vendored FLA Triton-AOT cubins kda_gate_cumsum → kkt → solve_tril → recompute_w_u → chunk_delta_h → chunk_gla_o) instead of the recurrence, exactly mirroring vLLM's prefill=chunk / decode=recurrent split (kimi_gdn_linear_attn.py:141). Requires VT_KIMI_DEVICE_KDA=1; decode (query_len==1) stays the recurrence. In the O(n²) recompute vehicle applying chunk to every step REGRESSED 122→102/128 (spec §18, #111) — the chunk is the vLLM-parity PREFILL half of the paged-incremental decode (spec §19: chunk-prefill ONCE then recurrent-decode the persistent state), where it carries vLLM's exact prompt-order arithmetic. Default OFF (parity-enabler) → the recurrence path is byte-identical |
VT_WHISPER_ENC_EAGER |
off (flash-tiled attention on) | Set to 1 to force the naive per-key block-reduction attention in the Voxtral/Whisper audio encoder instead of the default flash-tiled kernel. Rollback / A-B knob; token-identical to the default path |
VT_WHISPER_ENC_WARP |
off (flash-tiled attention on) | Set to 1 to force the warp-scoped online-softmax attention (vt::AttentionDenseFast, the pre-flash default) in the Voxtral/Whisper audio encoder. Rollback / A-B knob; bit-identical to the default flash-tiled kernel and ~1.15x slower than it |
VT_WHISPER_ENC_FA2 |
off (opt-in) | Set to 1 to run the Voxtral/Whisper encoder self-attention on the vendored FlashAttention-2 forward's TENSOR CORES (vt::AttentionDenseFa2) instead of the default scalar flash-tiled kernel. 5.50x faster encoder forward (731.7 -> 133.0 ms; 115.8x on the attention kernel alone, 19.28 -> 0.167 ms/layer), which takes the ENCODER FORWARD from 15.90x to 2.89x vLLM's whole TTFT (pinned oracle 555967922, TTFT median 46.02 ms). That is deliberately not stated as a TTFT ratio: our projector, merge and prefill are unmeasured, so they are missing from the numerator. NOT the default and NOT bit-identical: mma.sync reassociates the QK^T and PV reductions and FA-2 computes softmax as exp2f on a log2-scaled score, so results differ inside the bf16 envelope. Teacher-forced against the fixture's oracle it PASSES the ratified 0.5-nat near-tie band (worst 0.125 nats, 0 over-band) but with 3 divergent positions where the default has 0. Which of those differences flips the tokens is NOT established — the once-stated cause (bf16 P before the PV MMA) was refuted by mutation: forcing exactly that rounding into the scalar kernel left the tokens unchanged. Adoption is a developer decision — see .agents/specs/multimodal-speed.md §17.5 |
VT_FA2_DENSE |
on (when FlashAttention-2 is compiled) | Set to 0 to disable the vendored FA-2 fast path inside vt::AttentionDenseFa2, which then falls back to vt::AttentionDenseFlash. Same-binary A-B arm for the encoder attention; the op stays total either way |
VT_WHISPER_ENC_REMARSHAL |
off (encoder weights resident) | Set to 1 to disable device-resident encoder weights and re-marshal (host f32->bf16 convert + H2D upload) all Whisper/Voxtral encoder weights on EVERY forward, restoring the pre-residency behavior. Rollback / A-B knob; byte-identical output (moves data only). Default residency uploads each encoder weight once and reuses it, removing ~648 ms of per-call host marshalling from the encoder forward |
VT_QWEN3VL_ATTN_WARP |
off (flash-tiled attention on) | Set to 1 to force the warp-scoped online-softmax attention (vt::AttentionDenseFast, the pre-flash default) in the Qwen3-VL / Qwen3.6-27B vision tower per-frame self-attention instead of the default flash-tiled kernel (vt::AttentionDenseFlash). Rollback / A-B knob; the flash-tiled path is bit-identical to the warp path (tower output token-identical, goldens unchanged) |
VT_QWEN3VL_ATTN_EAGER |
off (flash-tiled attention on) | Set to 1 to force the naive per-key block-reduction attention (vt::Attention) in the Qwen3-VL / Qwen3.6-27B vision tower instead of the default flash-tiled kernel. Rollback / A-B knob; token-identical to the default path |
VT_DEVICE_KV_CACHE |
on (CUDA) | Host-side KV cache instead of the on-device one |
VT_GPU_SAMPLE |
on (CUDA) | Host-side sampling instead of on-GPU sampling |
VT_ROCM_Q8K_BLOCK |
unset (cooperative only on queue-device-resolved gfx1100) |
Exact 0 selects the permanent legacy A/B arm. Exact 1 forces the cooperative candidate diagnostically. Explicit 1 outside validated gfx1100 is not a support claim or a default claim. Unset keeps the legacy arm on gfx1200, gfx1201, unknown architectures, and architecture-resolution failure. Every other value is refused. See Q8_K activation quantization |
VT_ROCM_QUANT_WMMA |
on | ROCm keep-quant Q6_K and Q4_K prefill GEMM: the RDNA4 rocWMMA int8 tile arm (KQuantGemmKWmmaQ6K, KQuantGemmKWmmaQ4K), queue-device-resolved gfx1200/gfx1201 only and only when m >= 16 && n >= 16 (at least one full 16-wide tile in each dimension, not exact alignment — the scalar KQuantGemmK remaps its launched index to enumerate only the remainder the WMMA corner leaves untouched (the bottom strip past the aligned row boundary, then the right strip past the aligned column boundary), so a non-16-multiple M/N still takes the WMMA arm over its floor(M/16)xfloor(N/16) corner without paying a full-m*n-grid launch for the fill). 0 forces the scalar Dp4a arm this row's spec (KERNEL-QUANT-CIQ-GEMM-ROCM-RDNA4, issue #2109) is chasing a performance gap against, on every architecture and every shape. Bit-identical to the scalar arm by construction (the WMMA tile's raw int8 dot is scaled and reduced in the same integer arithmetic the scalar path uses; only the one f32 scale product per superblock is shared with it) — this is a same-binary A/B, not a correctness fallback |
VT_GEMV_MMVQ |
off | Exact 1 routes dense m == 1 Q4_K, Q5_K, and Q6_K ROCm keep-quant calls through the MMVQ GEMV arm. The arm folds activation quantization into the GEMV prologue when the output row count and scratch size fit the fold limits. Unset and every other value keep the baseline route. VT_GEMV_MMVQ_FOLD_MAX sets the inclusive output-row limit. Its default is 512; empty, malformed, nonpositive, and overflowing values use the default. Both values are read for each call, so graph capture records the selected route. |
VT_NORM_QUANT_FUSED |
off | Exact 1 lets a contiguous ROCm RmsNorm f32 or bf16 output produce Q8_K activation scratch for one immediately matching dense Q4_K, Q5_K, or Q6_K MatmulBTQuant call on the same queue and host thread. Unset and every other value keep standalone activation quantization. A mismatch consumes the one-shot token and keeps the standalone route. |
VT_ATTN_PREAMBLE_COOP |
off | =1 selects the warp-per-item cooperative attention preamble arm (AttnQkNormRopeGateCoopK) on ROCm, mapping one warp per token item instead of the donor walk (AttnQkNormRopeGateK); read once per process like the sibling arms |
VT_GDN_COLPERM_KEEP_QUANT |
off | =1 keeps the column-permuted ssm_out/out_proj tensor as Q5_K in tiled order (no ReorderVCols) and permutes the 4096-element GEMV input at runtime instead; the column reorder cuts across Q5_K block boundaries, so the weight cannot be permuted in place. Saves ~4x weight bandwidth (Q5_K ~5 MB vs bf16 20 MB per call) |
VT_GDN_ROWPERM_KEEP_QUANT |
off | =1 keeps the row-permuted V-head GDN projections (in the tiled order the row permutation produces) as K-quant instead of expanding to bf16 at load; the runtime gather supplies the permutation. Opt-in; the default reorders then expands |
VT_GDN_SCAN_COOP |
off | =1 selects the warp-per-row cooperative GDN scan arm (GdnScanCoopK) on ROCm, mapping one warp per output row instead of the donor walk (GdnScanK); read once per process like the sibling arms |
VT_GDN_PACKED_DECODE |
on (CUDA GDN) | Unpacked GDN decode path |
VT_GDN_DECODE_BV |
32 (CUDA GDN decode experiment) |
Exact 16 selects the byte-identical 16-value fused-recurrence tile; unset and every other spelling keep the 32-value schedule. Experimental opt-in; no release or cross-hardware default change |
VT_GDN_DECODE_SWIZZLE |
0 (CUDA GDN decode experiment) |
Exact 1 enables the shared-memory bank swizzle only for the BV=16, Dv=Dk=128, eight-lane production geometry; all other values and shapes keep the incumbent layout |
VT_GDN_DECODE_REGSTATE |
0 (CUDA GDN decode experiment) |
Exact 1 retains each lane's recurrence row in registers, only when the BV=16 shared swizzle is eligible. Other geometries and spellings use the shared-state implementation |
VT_GDN_OUT_BF16 |
on (CUDA GDN, EVERY arm since 5ae2c100f) |
Resolves the dtype of the two tensors the GDN block writes and the gated RMSNorm then consumes: dcore, the recurrence output, and z, the gate. =0 is the same-binary rollback that restores f32 for both. Until GDN-MOE-BF16-OUT (#1168) this variable only reached the DENSE arm, because GdnOutDType(bool dense_model) took a model-shape argument and all three call sites passed cfg.num_experts == 0, so a MoE checkpoint got f32 whatever the environment said. That argument is gone, the environment is now the whole decision, and the default therefore CHANGED for every MoE checkpoint. Mirrors vLLM, which keeps core_attn_out and the z gate bf16 on every arm. No speed number is established for either value. The spec's own `VT_GDN_OUT_BF16=0 |
VT_CONV_REG |
on (CUDA GDN) | The non-register-tiled short causal convolution |
VT_CONV_EXACT_CHUNKS |
on (CUDA GDN prefill) | Use =0 for the legacy sequence-serial causal-conv mapping; default mirrors vLLM's exact (sequence, 8-token chunk) descriptor and is byte-identical |
VT_CONV_CHANNEL_TILE |
0 (CUDA GDN prefill experiment) |
Selects the same-binary register-window causal-conv experiment: 0 is the sealed runtime-width one-channel baseline, 1 specializes the production width K=4 at the same 128-channel tile, and 2 specializes K=4 while assigning two channels per thread (256-channel tile, retained as the slower control). Only the exact values 1 and 2 opt in; unset, invalid values, and non-K=4 shapes fall back to 0. Experimental only; no default flip is implied by a microbenchmark win |
VT_GDN_POSTCONV_TOKEN_TILE |
0 (CUDA GDN prefill experiment) |
A non-0 value selects the byte-identical 16-token, per-head fused post-conv schedule. It remains opt-in pending the combined K4+post-conv Pareto measurement in issue #206 |
VT_MODELOPT_W4A4 |
0 (Qwen3.6 dense ModelOpt NVFP4) |
ModelOpt NVFP4 checkpoints ship a per-tensor input_scale next to every projection. Consuming it sets Nvfp4Weight::alpha, which flips IsTrueW4A4() and routes the weight to the fp4-ACTIVATION GEMM; on nvidia/Qwen3.6-27B-NVFP4 that produced incoherent text, so the default leaves alpha at 0 and takes the W4A16 weight-only dispatcher (verified coherent). Set 1 to consume input_scale and take the W4A4 path |
VT_LMHEAD_FP4 |
on (Qwen3.6 dense NVFP4 lm_head) |
Keeps a ModelOpt/compressed-tensors NVFP4 output head PACKED (Qwen3_5DenseWeights::lm_head_fp4) so the logits GEMM reads K*N/2 + K*N/16 bytes per step instead of the 2*K*N of a dequantized bf16 operand (~0.715 GB vs ~2.543 GB on nvidia/Qwen3.6-27B-NVFP4), and the operand keeps its on-disk [N,K] orientation instead of forcing the row-major NN GEMM that has no nvjet_sm121 kernel. Mirrors vLLM, which resolves a quantized lm_head through ModelOptNvFp4W4A16LinearMethod (modelopt.py:2491-2496,2508-2536) and never materializes bf16 (logits_processor.py:98-133). Measured peak host RSS 21.06 -> 19.36 GiB (-1.70 GiB: the bf16 head is 2,543,206,400 B = 2.368 GiB, the packed head 715,264,000 B = 0.666 GiB) — measured before #150 changed the RSS accounting, so the figure is owed a re-measurement. On a backend with NO fp4 GEMM the head additionally keeps ONE dequantized 2*K*N bf16 operand, built at prepare time; no other NVFP4 projection keeps one, so a quantized tower is never expanded (issue #203). The head is W4A16 under BOTH spellings: the on-disk activation divisor (input_scale / input_global_scale) is dropped for the head unless VT_MODELOPT_W4A4=1, because vLLM's ModelOptNvFp4W4A16LinearMethod deletes it (modelopt.py:1365). =0 is the same-binary rollback to dequantize-at-load. BF16, FP8, GGUF and tied heads are unaffected either way (row PERF-27B-LMHEAD-FP4, issue #213) |
VT_MOE_ROUTER_WARP |
on (CUDA, ungrouped MoE router, E in {32,64,128,256}) |
Runs the router top-k as ONE WARP per token with the whole logit row in registers — no shared memory, no __syncthreads(), and one global read of the row instead of two — where the incumbent 256-thread block kernel spends 37 barriers and 3 KiB of dynamic shared memory per token at E=256. BYTE-IDENTICAL, not approximately: the block reduction's s = 128, 64, 32 levels are all multiples of the warp width, so they never cross a lane and are exactly a per-lane halving tree over expert = lane + 32*slot, which makes the float additions, their operands and their association identical (derived per width in .agents/specs/moe-router-topk-single-warp.md §5, executed with no GPU by tests/vt/test_moe_router_warp_map.cpp). This is where the earlier "vLLM's topkGating reorders the softmax so it is off-limits" conclusion was too strong — that holds for vLLM's own CONTIGUOUS lane map, not for this one. =0 is the same-binary rollback to the block kernel. Grouped (noaux_tc) routing, the Serial byte-exact oracle, and every other E are untouched either way (issue #378) |
VT_FA2_PREFILL |
on (CUDA) | The portable prefill attention instead of the vendored FA2 |
VT_FA2_DECODE |
on (CUDA) | The portable decode attention instead of the vendored FA2 |
VT_FA2_DECODE_4B |
on (CUDA, Qwen3.5-4B) | The portable paged decode attention instead of the ratio-4 vendored FA2 path; the 27B and 35B selectors are unchanged |
VT_FA2_DECODE_GQA_SWAP |
on (CUDA, Qwen3-dense d128 decode) | Applies vLLM's FA2 seqlenq_ngroups_swapped decode optimization in the d128 varlen launcher (launches (batch, kv_heads) not (batch, hq), packing the ngroups query heads into seqlen_q); =0 opts out and restores the byte-identical plain-varlen reduction. Flipped default-on 2026-08-06 (KERNEL-FA2-GQA-SWAP-FLIP): SACRED Qwen3-0.6B/4B 16/16 token-identical, binding q3mxfp4 tput c2-c8 0.911/0.919/0.913→0.922/0.930/0.942 with no c1 regression, mem 2.614x less |
VT_FA2_SPEC_DECODE |
on (CUDA, spec-decode verify) | Keeps a classified uniform-qlen speculative VERIFY (PagedAttentionArgs::uniform_spec_query_len) on the d256 FA-2 split-KV DECODE lane instead of the num_splits=1 prefill ladder — spec-as-decode, SPEC-DFLASH2 W10 (#1857), mirroring FlashInfer's supports_spec_as_decode. =0 restores the prefill-lane routing in the same binary. Non-byte-exact vs the prefill route only when the split heuristic picks num_splits>1, because the split combine reorders f32 partials — the same near-tie class as VT_FA2_DECODE_GQA_SWAP |
VT_FA2_DFLASH_BLOCK |
on (DFlash/DFlash2/DSpark draft propose) | Routes the DFlash draft block's ATTENTION through the SHARED paged seam — vt::ReshapeAndCache writes the block's own (1+k) K/V into the draft store's pages, then one vt::PagedAttention reads [0, ctx+1+k) — instead of the bespoke vt::DFlashPagedBlockAttention op, whose block K/V live in no paged cache and therefore cannot reach any split-KV lane. SPEC-DFLASH2 W11 (#1890); the lane it leaves cost 449.7 us/call x 5.1 calls/step against SGLang's 14.3 us/call for the same work, while W10's verify on the FA-2 lane runs at 17.1. On CUDA it also admits the read onto the d128 FA-2 split-KV DECODE lane. =0 restores the bespoke op in the same binary. BYTE-IDENTICAL on CPU — the two kernels are the same three-pass online softmax in the same j-ascending order over the same bf16 bits, asserted element-for-element by test_qwen3_dflash_block_route across every mask the DFlash layer resolver produces. On CUDA it is the near-tie class this op already carries, for which VT_DFLASH_ATTN_BLOCK=1 is the bit-identical rollback. On a CUDA build where the FA-2 draft arm cannot engage the routed read is still correct but lands on the prefill-class kernel rather than the warp kernel it left, which is unmeasured; =0 is the rollback for that case too |
VT_CPU_MATMUL_STEAL |
off | Skip the decode-shape chunk-grid collapse in the CPU matmul so work stealing stays live (the collapse is a faithful ggml port that turns a decode GEMV into nth equal static chunks gated by the slowest core). Same-binary A/B; byte-identical either way, since which thread computes which output changes no reduction order. Not yet measured, so the default stays off |
VT_CPU_ELEM_KN_REPACK |
off | Transpose elementwise (f32/f16/bf16) [N,K] matmul weights to [K,N] at load, so MatmulBT takes the transpose-free CPU micro-kernels (1.16x to 1.30x, byte-identical). Off by default because the repacked bytes are transposed and only the CPU MatmulBTKernel honours the marker; enabling it for a non-CPU consumer would read garbage |
VT_CPU_REF |
off | Set on to force the portable reference path (dequantize-everything oracle), the standard "is this a kernel bug?" bisect switch |
VT_CPU_MATMUL_TIER |
auto |
Force the elementwise CPU GEMM tier in the same adaptive binary: portable/neon on arm64, or portable/sse2/sse2+f16c/avx2/avx512 on x86_64. An unsupported or cross-architecture tier fails closed after the exact OS capability probe |
VT_CPU_Q8_DOT |
auto |
Select the Arm Q8_0 x Q8_0 inner dot: portable, compiler-generated sdot, or the Linux-only a76-asm. auto selects the A76 schedule on Cortex-A76, otherwise SDOT on any Linux/Darwin DotProd-capable host, otherwise portable. A forced unsupported tier fails closed |
VT_CPU_PAGED_ATTN_NEON |
off | Set 1 to run the aarch64 NEON lane for the CPU paged-attention K dot-product and V accumulation (src/vt/cpu/cpu_paged_attn.cpp). Not bit-identical to the scalar path (the K reduction runs four fused-multiply lanes): outputs agree to f32 noise plus one bf16 output ulp, and greedy chains can flip near-ties — the kolibri gates adjudicate them; the decode-bench anchor is byte-exact and has no near-tie instrument yet, which is why the default is off (docs/bench-evidence/kolibri1-neon-paged-attn-20261008.md). Measured: attn_core halves (4.9→2.3 s at t32 over the 64-step decode bench). Widths that are not multiples of 4 in head or v-head dim, and every non-aarch64 build, use the scalar path regardless |
VT_CPU_QUANT_MMLA |
auto |
Select the Arm i8mm quant-dot tier. portable (also 0/off/false) disables it; i8mm (also 1/on/true) forces it and fails closed unless Linux HWCAP2_I8MM or Darwin hw.optional.arm.FEAT_I8MM is present |
VT_CPU_QUANT_REPACK |
auto |
Select the Arm i8mm+DotProd Q8_0 repack tier. portable (also 0/off/false) disables it; i8mm (also 1/on/true) forces it and fails closed unless the exact Linux HWCAP/HWCAP2 or Darwin sysctl capabilities are present |
VT_DFLASH_CTX_MAX_TOKENS |
unset (the resolved cap is max_model_len + 1 + k, itself capped at an 8 GiB budget shared across --max-num-seqs, so 8 GiB / max_num_seqs per request) |
Caps the DFlash/DFlash2/DSpark draft's CONTEXT STORE in tokens, PER REQUEST. The store is a per-request, per-draft-layer bf16 pool holding K and V, so its size is tokens x num_draft_layers x num_kv_heads x head_dim x 2 x 2 bytes per concurrent request, and the runner builds ONE PER BATCH ROW — so what the device holds is that number times --max-num-seqs, and gpu_memory_utilization accounts NONE of it. The default budget is therefore the AGGREGATE rather than a per-request allowance: a 262144-token --max-model-len would ask for about a gigabyte per request, and unbounded device residency has OOM-rebooted a box here (#1647). Set this BELOW the advertised context to bound that memory, or ABOVE the default share to speculate further into a long context — the aggregate arithmetic (tokens x bytes_per_token x max_num_seqs) is then yours. The effective limit AND the aggregate it costs are announced at startup either way. A request whose context passes the limit keeps running WITHOUT speculation, on the target alone — never refused, and never fatal: before #1919 the store was a fixed 4096 slots unrelated to --max-model-len, and the first prompt above it threw from inside an EngineCore step and left the server answering [request submitted to a stopped AsyncLLM] for every later request. Mirrors upstream, whose draft context lives in the engine's own paged KV cache sized cdiv(max_model_len, block_size) and whose speculators emit an EMPTY draft where they cannot serve a request |
VT_DFLASH_PAGED |
on (CUDA, DFlash spec-decode) | The materialized [context;block] draft forward instead of the fixed-capacity paged draft-KV store read through vt::DFlashPagedBlockAttention (bit-identical; only the DFlash single-request propose path) |
VT_SPEC_DECODE_GRAPH |
on (CUDA, spec-decode) | 0 runs the T=1+k speculative VERIFY eagerly instead of replaying a captured graph of that shape. Mirrors vLLM's uniform-decode dispatch, whose captured length is 1 + num_speculative_tokens; output is byte-identical either way, so this is a cost switch for A/B |
VT_DSPARK_DEVICE_SAMPLE |
on (DSpark spec-decode) | 0 restores the HOST sequential Markov loop. The device path keeps the per-step bias and argmax on the GPU and is token-identical to the host loop (same lowest-index tie-break, same base row, same d2t-mapped chain), so this is a cost switch for A/B only |
VT_SPEC_GRAPH_MAX_QLENS |
2 |
How many DISTINCT speculative query lengths one Qwen3.5 decode driver captures a graph for. Since ENG-CUDAGRAPH-BREAK W6 (#1374, #1020) the graph-eligibility predicate reads the step's ACTUAL uniform query length instead of the configured 1 + num_speculative_tokens, so a verify the scheduler clamped to a shorter draft prefix is captured at its own depth instead of running eager. That multiplies the ceiling on captured spec shapes by 1 + k, and each shape retains an [S, vocab] f32 logits block, so the widening is bounded. The default of 2 is the smallest value that admits anything new: the steady-state 1 + k plus one clamped length. 0 removes the bound; a larger value widens it. A step past the bound runs EAGER, which is what every clamped shape did before W6, and moves qlen_cap_declines |
VT_DFLASH_GRAPH |
on (CUDA, DFlash spec-decode) | The eager paged draft step instead of the captured/replayed draft-step CUDA graph (replayed==eager bit-identical; only the DFlash single-request propose path). Since ENG-CUDAGRAPH-BREAK W5 (#1335) the capture runs through the shared break-point seam, so VLLM_CPP_CUDAGRAPH=0 also turns it off; this name stays as the same-binary A/B lever for THIS driver alone |
VT_DFLASH_BOUNDS_DEVICE |
off (CUDA, DFlash spec-decode) | =1 reads the DEVICE seq_ext and the slot-map endpoints BACK from the device and asserts them against the host values the same call derived, refusing by name on a mismatch. It exists for the one defect class the pre-existing bounds checks cannot see: both are kCPU-guarded, so they confirm the host arithmetic was right and stay silent when the UPLOAD did not land on the tensor the paged read actually uses. OFF by default because the readback is a Download that SYNCHRONIZES, and this call sits on the no-sync path the route exists to keep — so it changes timing as well as checking, and it is a diagnostic run rather than a tuning knob. SPEC-DFLASH2 (#2274) |
VT_DFLASH_ATTN_BLOCK |
off (CUDA, DFlash spec-decode) | =1 selects the D12/D13 block-per-(query,head) draft paged-attention kernel instead of the D14 default warp-scoped online-softmax kernel (same f32-softmax math within the bf16 envelope; the D14 warp kernel is ~3x faster and closed the ~2% speed residual; only the DFlash single-request propose path) |
VT_DFLASH_ATTN_MMA |
on where the built arch set and the running device both have bf16 mma.sync (Ampere and later) |
=0 disables the bf16 tensor-core block-attention path and falls back to the CUDA-core form. The guard is both compile-time (__CUDA_ARCH_LIST__ all >= 800) and runtime (queried compute capability), so on a pre-Ampere board the path is already off and this flag changes nothing; it exists so the tensor-core and CUDA-core forms can be A/B'd on ONE binary, which is this project's benchmark protocol |
VT_DFLASH_ATTN_WARP |
off (CUDA) | =1 falls back to the older per-key warp-reduction block-attention kernel instead of the default chunked reduce-scatter form. Kept for the same-binary A/B that recorded the verdict |
VT_DFLASH_ATTN_KEYLANE |
off (CUDA) | =1 selects the one-key-per-lane block-attention form. MEASURED NEGATIVE and not a tuning knob: 28.90 s/step against the per-key warp kernel's 18.73 on the same binary (sm_110, MiniMax-H3 512x512/33f, seq 3224), 54% slower, because giving each lane a whole K row makes every K load 32-way scattered. Kept only because it is the experiment that located the real constraint |
| VT_ADOPT_DEVICE_BYTES | on (acts only where the backend advertises host-addressable device memory: Vulkan always, Metal and ROCm where the part reports unified memory. MEASURED on Vulkan only) | After a weight is uploaded, re-point its host buffer AT the device allocation instead of keeping a second copy. On a unified box the two copies come out of the same RAM. MEASURED on GB10, Qwen3.6-27B bf16 (50.89 GiB): with the mirror the process reaches VmRSS 100.759 GiB and drives the machine to MemAvailable 13.85 / MemFree 1.13 GiB of 119.6 GiB before it has even finished allocating; without it the same load completes at VmHWM 53.413 GiB. Qwen3-4B: 16.392 -> 9.607 GiB. Vulkan allocation is byte-identical either way. 0 is the same-binary A/B back to the two-copy behaviour, and it moves ONLY this lever: the VT_LOAD_DIRECT_UPLOAD post-upload source-page release still runs at 0 and on a backend without host-addressable device memory, so the two knobs stay independent. It is an adoption, not a release — the bytes survive at the device address and every reader sees them — so tokens are unchanged either way (test_opt_paged_engine on Vulkan is 6/6 token-exact, 96/96, both arms). REACH AND MEASUREMENT ARE NOT THE SAME SET, and this row used to read as if they were. The lever is gated on vt::Backend::DeviceMemoryIsHostAddressable(), read at both AdoptDeviceBytesAsHost branches in src/vllm/model_executor/models/qwen3_5_weights.cpp. That predicate defaults to false (include/vt/backend.h), so a backend has to opt in, and THREE now do. VulkanBackend (src/vt/vulkan/vulkan_backend.cpp) returns true unconditionally, because every allocation it makes is HOST_VISIBLE|HOST_COHERENT and persistently mapped. Since cffe59b02 (#1477, which moved the portable reference tier off the WIDER UnifiedMemory()) MetalBackend (src/vt/metal/metal_backend.mm) answers MetalContext::unified_memory(), i.e. dev.hasUnifiedMemory, and RocmBackend (src/vt/rocm/rocm_backend.hip) answers its own unified_memory_, i.e. a managed allocator or an integrated part with PageableMemoryAccess. So the lever ACTS on Apple silicon and on an integrated ROCm part, and nobody has measured it on either — every number above is GB10 through Vulkan. Read those two as reachable-and-unmeasured rather than as covered by this row's numbers (#1502). It is still inert on CUDA and CPU, neither of which overrides the default. That is now CHECKED rather than merely observed (#1635), and in two halves, because no CI job has a GPU on which a live CudaBackend could answer: a static_assert beside the class in src/vt/cuda/cuda_backend.cu fails the cuda-fat-build compile if CudaBackend ever declares its own DeviceMemoryIsHostAddressable, so its answer IS the inherited one; and a case in tests/vt/test_backend.cpp pins that inherited default to false on a subclass that declares no override, on every host lane. Do NOT cite tests/vllm/platforms/test_platform.cpp for this: its FakeUnifiedAddressablePlatform reports device_type() == kCUDA but its backend() returns vt::GetBackend(DeviceType::kCPU), so its CHECK_FALSE(gb10.backend().DeviceMemoryIsHostAddressable()) reads the CPU backend. tests/vt/test_cuda_backend.cpp observes the real backend but skips wherever no GPU is present, so it asserts nothing in CI. The CPU backend reports UnifiedMemory() == true while this narrower predicate stays false — the two properties disagreeing is the whole reason the second one exists |
| VT_QWEN35_ALIAS_HOST_WEIGHTS | on (Qwen3.5/3.6 family; acts only where the platform advertises host_memory_is_device_addressable() — a GB10-class integrated part today) | Hand a dense weight's HOST bytes straight to the device kernel instead of allocating a device copy and uploading into it. On a part whose kernels can dereference host storage that copy buys nothing and costs a second full resident copy of every dense weight out of the same RAM. MEASURED on GB10, Qwen3.8-2.4T-A95B UD-Q1_0 (369.97 GiB) on --device cuda: with the second copy the load completes at 61.20 GiB resident and then exhausts the 119.631 GiB box inside its FIRST forward, zero decode steps over seven attempts (issue #1299); with the alias the same run reaches 32/32 decode steps at peak RSS 97.75 GiB. The instrument counts 60.793 GiB of dense weight aliased rather than duplicated (first-forward totals, at call 1361, where re-homing plateaus) against ~9.2 GiB of misaligned GGUF borrows that decline and still stage. 0 is the same-binary A/B back to the staging behaviour, and it is not only a bisect lever: src/vllm/model_executor/models/laguna.cpp records a MEASURED GB10 penalty for reading system-allocated memory from the GPU rather than a cudaMalloc allocation, worst on a long-K low-parallelism GEMV, so a decode regression has to be separable from the workload. The substitution is otherwise indistinguishable — the aliased pointer is re-homed to 256 bytes, cuBLASLt's own CUBLASLT_MATMUL_PREF_MIN_ALIGNMENT_A_BYTES default, which dominates every explicit pointer gate in the CUDA kernels (the strictest asks 32). A DISCRETE device answers the predicate false and is byte-for-byte unchanged. No effect on CPU, Vulkan, Metal or XPU, whose platforms do not advertise the property |
| VT_QWEN35_STAGE_RESERVE_BYTES | 12 GiB (Qwen3.5/3.6 family; read once, and only by the SAFETENSORS loader) | The device memory a staged model has to leave behind for the KV cache and activations. The retag above is measurably slower than a true device copy on GB10, so a model that FITS should stage; #1299 recorded Qwen3.8-2.4T-A95B exhausting a 119.631 GiB box precisely BECAUSE the CUDA arm paid for its weights twice, so a model that does not fit must not. The rule is 2 * model_weight_bytes + VT_QWEN35_STAGE_RESERVE_BYTES <= device_total_bytes, and the factor of TWO is the substance: on a platform that does not release the host mirror after upload, a staged weight is ADDITIVE. MEASURED on the committed gate (scripts/dflash2-speed-gate.sh, Qwen3.8-27B bf16 + DFlash2 k=7, c=1, one boot, same oracle wheel): retagged 12.361 tok/s -> staged 15.029, against vLLM's 16.363, i.e. ratio 0.7587 -> 0.9185, +21.6%. IT IS DECIDED ONCE, AT LOAD, AND THAT IS THE WHOLE POINT. The first version of this policy asked Backend::DeviceMemoryInfo PER WEIGHT and compared LIVE FREE against a fraction of total; free falls as the host mirror loads, so the floor degraded into "stage the first N GiB, then stop" and captured 8.1% of the available 21.6% (a22030924, gate 13.3625). The rule was already written next door: include/vllm/platforms/interface.h says the budget is "TOTAL rather than FREE, because free at load time carries the page cache and whatever else the box is doing, which would make a load-time verdict a function of contention." Only the safetensors loader latches it, so a GGUF checkpoint never stages and keeps exactly the behaviour #1299 shipped — which also makes the file total a sound proxy, since a bf16 safetensors weight occupies its file size in host RAM while a GGUF one expands on load. Parsed with atoll: unset, empty, unparsable and <= 0 all fall back to 12 GiB rather than refusing, so a typo cannot silently shrink the reserve and double a model's residency. An unknown device total or an unknown model size answers NO. VT_QWEN35_ALIAS_HOST_WEIGHTS=1 pins the retag on regardless; =0 forces staging regardless. LOWER admits bigger models to staging; HIGHER is more conservative |
| VT_LOAD_DIRECT_UPLOAD | on | Load a weight the device consumes VERBATIM by VIEWING the safetensors mmap (OwnedBytes::Borrow, keep-alive on the mapping) instead of copying it into an owned host buffer first, so the device upload reads the file mapping and the load moves those bytes ONCE rather than twice. Only whole-range same-size copies qualify — a transpose, a dtype conversion, a dequant, a concatenation or a load-time repack always takes the copy path, and the helper re-checks numel * sizeof(dtype) == span and fails closed to the copy on any mismatch. 0 is the same-binary A/B back to copy-then-upload. Bytes are identical either way, so tokens are unchanged. MEASURED on GB10, Qwen3.6-27B bf16 (50.098 GiB), Vulkan, same binary both arms: the weight-load phase goes 19.27 -> 12.48 s warm (1.54x) and 52.62 -> 32.75 s cold (1.61x), load-and-one-token 30.39 -> 22.47 s warm and 62.98 -> 55.60 s cold, with every ON leg beating every OFF leg. Total bytes MOVED 100.196 -> 81.260 GiB: the host materialization pass drops 50.098 -> 31.162 GiB while the 50.098 GiB device upload is unchanged (the model still has to be uploaded once). 37.8% of this checkpoint qualifies; the rest is merged (qkv, gate_up) or transposed (lm_head) at load and correctly still copies |
| VT_LOAD_STATS | off | =1 prints one line per load phase (mmap+header, weights) with its wall time, plus the bytes the load MOVED: host_copy (source bytes materialized into an owned host buffer), borrowed (source bytes viewed in place by the direct-upload path) and device_upload (bytes copied host to device). Diagnostic only; it changes no numerics. Issue #150 |
| VT_VULKAN_ALLOC_STATS | off | =1 prints a device-memory line on every 1 GiB high-water crossing and a summary at exit: live buffer count, bytes REQUESTED by the caller, bytes COMMITTED by the driver (VkMemoryRequirements::size), peak live bytes, and the process/system context (VmRSS, VmHWM, MemAvailable, Cached) read from /proc. On a unified-memory device the Vulkan heap IS system RAM, so separating "the backend allocated it", "the process allocated it some other way" and "it is page cache" is the whole of a memory attribution. Diagnostic only; it changes no numerics. The counters themselves are always maintained (one relaxed atomic per allocation) and are readable from a test through vt::vulkan::DeviceAllocStatsSnapshot(). Vulkan-only |
| VT_VULKAN_DISPATCH_STATS | off | =1 traces every Vulkan compute submit to stderr (index, shader, workgroup count) BEFORE its fence wait, prints any wait over 200 ms, reports a running dispatch rate every 100 submits, and dumps a per-shader histogram at exit. Printing before the wait is what makes a HANG visible: a post-wait print never runs if the fence never signals, so the last line names the dispatch that hung. This is how the coopmat out-of-bounds load was found. Diagnostic only; it changes no numerics. Vulkan-only |
| VT_VULKAN_GEMV_UNROLL | 4 | =1 forces the un-unrolled decode GEMV body. Four independent accumulators keep four reads per lane in flight instead of one -- memory-level parallelism, not instruction count. It rides a specialization constant, so both arms are the same committed module and A/B in one binary. MEASURED 1.055x, 7 of 8 interleaved pairs. Worth noting it measured 5/8 and was REVERTED earlier the same day: that test ran while the GPU was only 26% busy, where a 10% GEMV win moves e2e by 1.4% and is unresolvable against this box's noise. After the ring fix made the run GPU-bound the same code reads 7/8. A negative result is regime-dependent. Vulkan-only |
| VT_VULKAN_GEMV_PACK | 2 | Load width for a 16-bit decode GEMV operand: 0 one element per load (2 B), 1 two through the buffer's 32-bit view (4 B), 2 four through a 64-bit view (8 B). Same bytes, same coalescing (32 lanes still cover 128 or 256 CONSECUTIVE bytes), so this cannot reduce DRAM traffic; it reduces LOAD INSTRUCTIONS. MEASURED, and the two regimes disagree by 7x, which is the point. In benchmarks/vulkan_gemv_ab.cpp over the 27B's own decode shapes (9 arms x 4 rotated passes, each arm paired against the width-0 baseline measured IN THE SAME PASS, because the box drifted 15.5% peak-to-peak between passes) width 2 reads 1.086x. In REAL 27B decode it reads 1.012x and 1.020x, two GPU-timestamp two-length diffs (output-len 36 minus 4 over 32 tokens), moving vt_matmul_vec from 90.0% to 91.1% of the 273 GB/s roof. The sweep re-reads one 356 MB buffer 320 times, so its DRAM rows and TLB stay hot and instruction issue is visible; decode streams 50 GB once per token, where DRAM is the whole story. e2e: 4.126 -> 4.157 tok/s, 5 of 6 clean pairs, against a 0.88% clean-leg noise floor. Degrades a width at a time when K or an operand byte offset is not aligned to it, and declines entirely for an f32 operand, which is already one element per 32-bit word. Unlike the row count this axis DOES change the answer's low bits, because it repartitions K across lanes; the opt-125m STRICT gate (6/6 token-exact) is what clears it. Vulkan-only |
| VT_VULKAN_GEMV_ROWS | 1 | Output elements per decode-GEMV workgroup, 1, 2 or 4; llama.cpp's NUM_ROWS (mul_mat_vec_base.glsl:90, set per vendor by rm_stdq in ggml-vulkan.cpp:4705-4719). Raising it lets one loaded activation element feed several weight rows. It cannot reduce DRAM traffic either -- the weights are read exactly once per token at any row count -- and the activation re-reads it saves were L2 hits. MEASURED NEGATIVE on GB10 and shipped OFF: 0.966x at 2 rows and 0.968x at 4, in all four passes and at every load width, most likely because halving the workgroup count halves the independent sequential read streams the memory controller sees while each surviving workgroup interleaves rows K*2 bytes apart. Kept as an axis only because llama.cpp raises it exactly on the AMD GCN and Intel parts this backend's portability case is aimed at, and deleting it would mean rediscovering it there. Bit-identical to 1 at any value -- each lane accumulates the same strided subset of K in the same order -- so it is a pure scheduling knob. Vulkan-only |
| VT_VULKAN_RING | 128 | Descriptor-set ring depth per pipeline, i.e. how many times one shader may be recorded before the batch must flush. This is what actually caps batch length: at 16 it forced a submit every 40-46 dispatches, because vt_rms_norm runs 112 times per forward pass and exhausted its ring seven times. At 128 a whole forward pass (368 dispatches) fits in ONE submit. Each avoided flush is a vkQueueSubmit plus a blocking vkWaitForFences, and a host profile puts 62% of on-CPU time in the kernel and NVIDIA driver against 14% in our own code -- so submits are the host cost. Lower it to bisect or to cut descriptor-pool size. Vulkan-only |
| VT_SERVER_VERBOSE | off | =1 turns on the server's chatty progress output, and is the umbrella switch the individual knobs fall back to: VT_SERVER_PREFILL_PROGRESS reads it when unset, so VT_SERVER_VERBOSE=1 alone enables chunked-prefill progress (rate-limited to ~2 Hz per request). examples/server sets it from its own verbosity flag, so it is user-facing rather than a kernel tuning switch |
| VT_ENGINE_STEP_LOG | off | =1 prints an EngineCore step heartbeat to stderr (unfinished count, whether the model executed, output rows, elapsed). Also on under VT_SERVER_VERBOSE=1. It is how you tell an idle engine from a stalled one: a run of model_executed=0 with a non-zero unfinished count is work the scheduler cannot admit. Diagnostic only; it changes no generation |
| VT_VULKAN_INFLIGHT | 2 | How many command-buffer submissions may be in flight at once. At 1 a flush submits and BLOCKS on its fence, so all host recording is serialized against an idle GPU; at 2 the flush returns and the next token records while the previous one runs, worth a MEASURED -1.41 ms/token on 27B (5/5 paired legs). Each slot owns a disjoint descriptor-ring slice, so raising this raises descriptor memory |
| VT_VULKAN_HOST_PROFILE | off | =1 prints a per-phase HOST time breakdown at exit (submit, record, descriptor update, bookkeeping, pipeline lookup, and time blocked on the GPU). Diagnostic companion to VT_VULKAN_DISPATCH_STATS, which covers the DEVICE side; use it to separate host cost from GPU wait |
| VT_VULKAN_BATCH | on | =0 forces the per-dispatch submit-and-wait path instead of recording many dispatches into ONE command buffer (with a pipeline barrier between them) and submitting once. MEASURED on GB10: decode 2.62x, faster in 8 of 8 interleaved pairs (24.8 -> 64.5 tok/s), 40-46 dispatches per submit, opt-125m still 6/6 token-exact. Batching is sound only because all three host-read paths drain the batch first: Backend::Copy and Memset (host memcpy over the mapped allocation), Synchronize, and Backend::FlushPending, which op_provider.cpp calls before running a PORTABLE REFERENCE TIER kernel directly over device memory. Without that third one a host kernel would read stale bytes silently. It exists for the same-binary A/B and as the bisect lever. Vulkan-only |
| VT_VULKAN_SMART_BARRIERS | off | =1 records a vkCmdPipelineBarrier before a batched dispatch ONLY where a real dependency exists, instead of before every one. The batched path otherwise emits ~900 full COMPUTE->COMPUTE memory barriers per 27B decode token; llama.cpp instead calls ggml_vk_sync_buffers at 42 explicit call sites and tracks per-scratch-buffer need_sync flags (ggml-vulkan.cpp:3193, :8687-8748 @ pin 237ad9b96). This is that idea made general: each dispatch's READ set and WRITE set come from the committed SPIR-V's NonWritable decorations (vulkan_spirv.h § writable_mask), which glslang enforces, and a barrier is skipped only when none of read-after-write, write-after-write or write-after-read holds against everything recorded since the previous barrier. It DEFAULTS TO EMITTING: buffers are compared whole rather than by byte range, an operand that is both read and written counts as a write, and an unrecognised binding counts as writable, so every uncertainty costs a barrier rather than correctness. Both arms live in one binary because a cross-BUILD comparison of two barrier policies has already produced a false 1.2x in this campaign -- and because a policy that drops a real dependency computes WRONG NUMBERS, which is invisible on llvmpipe and was invisible for a whole benchmark block the last time it happened here. Vulkan-only |
| VT_VULKAN_GEMV | on | =0 forces the decode GEMV onto the portable scalar GEMM instead of vt_matmul_vec. The vec tactic gives each output element a WORKGROUP whose lanes stride K, so adjacent lanes read adjacent addresses; the scalar kernel puts one invocation on each output element and loops K there, which for the MatmulBT weight layout leaves adjacent lanes K*2 bytes apart, each pulling its own cache line for 2 bytes of it. Selected only for MatmulBT at M=1 with K at least one workgroup wide -- the other orientation is already coalesced and the vec shape would make it worse. MEASURED on GB10: GEMM ms/call improves in 7 of 8 interleaved pairs, median 1.8x. Its K reduction is a tree, so unlike the scalar kernel it does NOT share the CPU's accumulation order; the opt-125m STRICT gate (6/6 token-exact) is what clears it. Vulkan-only |
| VT_VULKAN_RMSNORM | auto | Which vt::RmsNorm SPIR-V module runs: wide forces the 1024-invocation subgroup-reducing one, base forces the portable 128-invocation one, unset lets the device capability decide (1024 invocations on the X axis plus compute subgroup BASIC and ARITHMETIC). The wide module exists because RmsNorm dispatches ONE WORKGROUP PER ROW and a batch-1 decode step has exactly one row: on Qwen3.6-27B that put 128 invocations on a 5120-wide row, four warps of one SM, with the rest of the GPU idle. MEASURED on GB10 by the two-length GPU-timestamp diff: 0.0611 -> 0.0123 ms/call, 7.88 -> 1.59 ms/token, and paired decode 241.9 -> 235.6 ms median TPOT. The tell that it was OCCUPANCY and not the reduction is that the SAME shader costs 0.066 ms/call during PREFILL, where 32 rows give it 32 workgroups and 32x the data. It exists for the same-binary A/B and so the unit gate can exercise the fallback on hardware that would always pick the wide arm. Vulkan-only |
| VT_VULKAN_MATMUL_NCOLS | 4 | Output columns each lane of the portable scalar GEMM computes, in the [K,N] (non-transposed) orientation only. At 1 the kernel is the flat one-invocation-per-output-element body; above 1 a workgroup takes 128*NCOLS CONSECUTIVE output columns of one row, so at each step of K it reads a contiguous run of that many elements instead of 128. This is the ONE decode GEMM that cannot reach the vt_matmul_vec tactic, because in [K,N] the lanes are already coalesced and the GEMV shape would make them strided; on the 27B it is the lm_head, m=1 k=5120 n=248320, 2.54 GB moved per token. MEASURED on GB10, 27B decode, ms/call medians over interleaved replicates: NCOLS 1 = 12.48, 2 = 12.46, 4 = 11.54, 8 = 12.81, with 4 winning 6 of 6 interleaved pairs against 1. Blocking is a TRADE, not a monotone win: at 8 the dispatch falls to 243 workgroups (~31k threads) and the device runs out of work to hide memory latency with faster than the longer contiguous run buys back. It rides a specialization constant, so every arm is the same committed module and they A/B in one binary. Every arm is BIT-IDENTICAL -- each accumulator owns one output element and sums the whole K sequentially, which is the CPU kernel's order -- so this kernel keeps the byte-exact tier that the coopmat and GEMV tactics gave up; a memcmp gates that. Vulkan-only |
| VT_VULKAN_COOPMAT | on | =0 forces the Vulkan GEMM onto the portable SCALAR kernel instead of the cooperative-matrix (tensor-core) tactic. The coopmat path is selected only where the device reports the exact 16x16x16 bf16/bf16/f32/f32 SUBGROUP configuration, subgroup size is 32, both operands are bf16, and M, N and K are all multiples of 16. The whole-tile requirement on M and N is not a tuning choice: coopMatLoad reads a full 16x16 tile with no masking, so a partial tile reads past the operand and can fault the GPU. Ragged shapes fall back to the scalar kernel; this switch bypasses that selection entirely. It exists for the same-binary A/B in examples/vulkan-gemm-ab (measured 11.1x-32.9x on NVIDIA Thor) and as the bisect lever if a coopmat result is ever suspect. Vulkan-only |
| VT_GLM5_NEXT_DEVICE_EXPERTS | off (opt-in) | =1 lets Glm5NextForConditionalGeneration (GLM-5.3-Flash) accept a non-CPU queue and route its routed-expert keep-quant GEMM to the device, against banks made resident by dense_attn::ResidentWeight. IT IS OFF BECAUSE THE PATH IT ENABLES IS MEASURED TO CRASH, not because it is unmeasured. On dgx:gpu0 against the published 101.24 GiB UD-Q2_K_XL artifact, ALL THREE --device cuda legs died with SIGSEGV (rc=139) having emitted no token, interleaved against three --device cpu legs that all emitted Paris. from the same binary. The mixed-residency reading of those legs is FALSIFIED and the cause is elsewhere: the two log lines that suggested it are once-flags, and the process dies in StoreCaches, which host-stores into the runner's cudaMalloc KV pages after the forward has already returned. That defect is older than this knob and only became reachable when the non-CPU refusal above it was removed; see .agents/specs/glm5-next-flash.md O49 and #2480, which owns the fix. The default is the refusal the tree carried before the arm existed, because turning a clean named error into a segfault is strictly worse for a user. Set this only to debug that crash; it is not a serving knob. Parsed strictly (1 and nothing else, not the usual first-character rule) precisely because it opts into a crashing path. Inert on every other model and on --device cpu. See .agents/specs/glm5-next-flash.md O46 and #2464 |
| VT_GLM5_NEXT_DEVICE | off (opt-in) | =1 routes the entire GLM-5.3-Flash forward through Glm5NextDeviceForward, which dispatches embedding, RMSNorm, MoE combine, the k-pool indexer and lm_head through vt::* device ops on the queue, keeping MLA attention, mHC sites and the dense MLP as host-fallback islands (the kimi_linear_device.cpp single-queue pattern). This is a superset of VT_GLM5_NEXT_DEVICE_EXPERTS: when on, the whole forward delegates and the per-arm experts split is not reached. On a CPU queue the vt::* kernels use float32 accumulation where the host reference uses double, so the output agrees within a float-vs-double envelope rather than byte-exact. Inert on every other model. See .agents/specs/glm5-next-flash.md W9c-3 and #3175 |
Read-only observability; none change output.
| Variable | Default | What it does |
|---|---|---|
VT_DFLASH_GRAPH_STATS |
unset | Print DFlash draft-step CUDA-graph capture/replay counts to stderr |
VT_GRAPH_STATS |
unset (off) | =N prints the graph-dispatch and DFlash-route counters to stderr every N steps, from the shared step path in GPUModelRunner::execute_model. Two lines: [graph-dispatch] steps= uniform= uniform_spec= clamped_spec= ragged= ragged_mixed= ragged_prefill= ragged_spec= ragged_pct= spec_as_decode= capture_shapes= qlen_cap_declines= and [dflash-route] paged_seam= block_kernel= combined= last_combined_q= last_combined_k=. Both counter families existed before #2112 and neither had a reader outside tests/, which made E6 of .agents/specs/dflash2-batch-propose.md unrunnable as written and left the ragged-step share invisible on a server. ragged_mixed is the field #2117 turns on: it counts steps that carried a prefill row and a decode row together, which is the population that loses both its CUDA graph and its decode attention lane, and it is separated from ragged_prefill (ordinary prefill, nothing degraded) and ragged_spec (uneven verify widths, #1943's shape) because a single ragged count is consistent with all three and therefore distinguishes none. The three PARTITION ragged. qlen_cap_declines is in the line so that the free VT_SPEC_GRAPH_MAX_QLENS=0 A/B is interpretable: no throughput movement with declines in the default arm means something different from no movement with zero declines. Read-only; unset is byte-identical and costs one integer compare per step |
VT_OP_PROVIDER_STATS |
off | Print per-op provider (which backend served each op) statistics |
VT_TT_TRACE_DEBUG |
unset | =1 prints the Tenstorrent capture bisection traces to stderr: op entries ([TT-OP]), host readbacks (to_vector/EnsureHostBytes), device->device copies/zero-fills, and rope cos/sin cache lookups — all gated to fire only while a mesh-trace capture is active. Read-only diagnostics for the host-free decode investigation; byte-identical output when unset |
VT_TT_GDN_DECODE |
unset | =chunked runs kGdnDecode as one T=1 chunk_gated_delta_rule call instead of the composed rank-1 matmul+eltwise step. Diagnostic arm of the W2 composition measurement (spec tenstorrent-gdn.md): the fused op pads T=1 to its 64-token chunk, measuring 3.164 vs 1.139 ms/step at B=8 GQA 2:8 Dk=Dv=128, so composed is the default. Same numerics envelope either way |
VT_TT_GDN_STATE_PROBE |
unset | =1 prints one [GDN-STATE-PROBE] line per EAGER kGdnDecode call to stderr: the step counter, the cache slots, the batch-heads, the gathered initial state's element count, an FNV-1a checksum of its bytes and the absolute-sum. The capture-warmup redesign's instrumented state probe (spec tenstorrent-capture-warmup-redesign.md, W2): a frozen GDN recurrence prints ONE identical checksum at every step of a sequence, an advancing recurrence prints changing ones. Read-only diagnostics; the download it needs is the read a trace capture refuses and replays run no host code, so capture steps print nothing, and an unset run pays nothing |
VT_TT_GDN_REPLAY_TAP |
unset | =1 prints the GDN region's replay attribution to stderr, per Qwen3_5DenseDecodeGraph step (cold / capture / replay): one [GDN-TAP] line per layer carrying the residual stream's post-layer FNV-1a checksum, min/max and zero count (the capture-safe in-region aux-tap buffer, so a REPLAY's taps are the replayed values), the logits' argmax/zero-count/first-five, and one [GDN-SHADOW-PROBE] line per GDN state cache (ssm and conv) checksumming the CURRENT device shadow BETWEEN steps — a slot frozen at the capture-time values prints one identical checksum at every replay (the stale-binding signature), an advancing slot prints changing ones. The region-replay row's instrument (spec tenstorrent-gdn-region-replay.md): the FIRST layer whose tap reads zero at a replay names the defect's layer. Read-only diagnostics: every read is a device→host download taken between steps and never during capture, and an unset run pays one getenv |
VT_TT_STEP_PHASES |
unset | =1 prints one [TT-STEP-PHASE] line per Qwen3_5DenseDecodeGraph step to stderr — the 27B captured step's host phase clock (spec tenstorrent-27b-step-decompose.md): the per-step warmup-refresh passes and their wall time, embed, the capture pass split into begin/body/end (the tt-metal trace build), the replay launch, the step wall, the free-DRAM read-out (the ~950 MB/request retention staircase, tt-metal#57970), and the gap to the previous step; the seam-side sync lines name the completion wait (the blocking read that pays the replay's device tail). =sync additionally brackets each of the 64 layers in DenseForwardLayers with a blocking queue drain, so each layer line carries its true DEVICE time — a sampling leg (serialized), not a wall measurement; it is inert while a trace capture is open. Read-only either way: the anchor's token stream is byte-identical with the knob set, and nothing runs when it is unset |
VT_OP_PROVIDER_DISABLE |
(none) | Comma-separated provider names to disable, forcing fallback (diagnostic) |
VT_SERVER_PREFILL_PROGRESS |
off | =1 prints chunked-prefill progress to stderr, rate-limited to roughly 2 Hz per request. =0 explicitly disables it even when VT_SERVER_VERBOSE=1 |
VT_GLM5_DIAG |
off | =1 traces one GLM-5.3-Flash forward step to stderr: the step shape and its token ids, the embedding gather, every bridged layer's per-tensor norms, ONE line per decoder layer, the collapsed hidden state entering the head, the first lm_head chunk, and the logits with their top-5 (id, value) pairs and the top1-top2 MARGIN. =2 adds the sublayer trace inside each layer (the mHC collapse, both norms, the attention arm's output, the fold, the feed-forward output). Each line carries NaN, Inf and zero counts, min, max, mean, stddev, L2 and the first four values, because the four defects that all emit token id 0 -- NaN logits, finite uniform logits, finite wrong logits, and a broken sampler -- are separated by VALUES and not by a verdict. The margin is printed because a discrete selection has bimodal error: a top-1 that wins by 0 is a tie the argmax broke by index and is a different fact from one that wins by 3. Read-only; every call site is behind the level check, so an unset run is byte-identical, which is also why level 1 and above print an INSTRUMENT ACTIVE banner naming the raw environment string -- a probe that silently did nothing would otherwise read as a clean tensor. Issue #2241 |
VT_NEMOTRON_H_DIAG |
off | =1 prints one line per step and per layer of NemotronHPagedForward to stderr: the decode/prefill split, the recurrent slot indices and the has-initial mask, then per Mamba2 layer the L2 norm of the state gathered IN and the state written OUT, and per layer the last row's mixer and residual norms. It is what separated a broken recurrent carry from a stale input id on this model (#1157): the carry read identical on host and GB10 while layer 0's embedding row did not. Every download it does is inside the guard, so an unset run pays nothing |
VT_NEMOTRON_H_DEVICE_MOE |
on | =0 routes NemotronH's NVFP4 MoE layers back through the host reference instead of the device Marlin arm. A bisect lever for a device-vs-host divergence on this model, not a configuration: it trades the whole MoE tower's throughput for the arm the CPU gate proves token-exact |
VT_GDN_VALIDATE |
off | Run the GDN validation/cross-check path (slower; for kernel debugging) |
VT_FP4_AUTOTUNE_VERBOSE |
off | Log the NVFP4 GEMM autotuner's tactic selection |
VT_H3_PROGRESS |
unset | Trace the MiniMax-H3 denoise loop's phases to stderr: which forward path was taken (device vs the CPU reference), how long the ONE-TIME device weight staging took, and per-step forward seconds with the sequence length. A real-checkpoint run spends its minutes in exactly one of those phases, and this says which without guessing — it was added after GPU-utilization counters proved unreliable on Tegra-class boards |
VT_POOL_BYPASS |
off | =1 makes every device-scratch pool allocation an exact-size driver Alloc and every release a real Free, so compute-sanitizer can see tensor boundaries and use-after-free that the caching, size-class-rounding pool hides. DEBUGGING ONLY: it reinstates the per-op cudaMalloc/cudaFree device-sync storm the pool exists to remove, so it is never a timing configuration |
VT_POOL_BORROW |
on | =0 restores the pre-#1922 exact-class-only reuse in the device-scratch pool, where a freed block could only ever serve a request in its OWN size class. ON by default and there is no supported reason to turn it off: with the borrow off, what the pool retains is a function of how many distinct shapes the traffic has shown rather than of how much one step concurrently needs, and nothing on the serving path ever gives any of it back. Twelve sequential requests through the production engine entry, with the LARGEST one first so every later request demanded strictly less than one already served, retained 560 243 B after the first and 2 867 707 B after the twelfth with the borrow off, against 624 899 B with it on. This exists so the A/B that MEASURES what the borrow is worth runs on ONE binary, which is what a same-binary A/B requires, and so tests/vllm/v1/worker/test_engine_scratch_steady_state.cpp has an in-tree red arm. Not a configuration |
VLLM_LTX2_POOL_DRAIN |
on | =0 suppresses the LTX-2.5 render's per-phase device scratch-pool drain. The drain is on by default and there is no supported reason to turn it off: a denoise phase leaves the pool holding every activation size class it touched, the next phase (twice the resolution) and the VAE decode allocate different classes, and on an uncapped pool those blocks are headroom the next stage never gets. This exists so the A/B that MEASURES what the drain is worth runs on ONE binary, which is what a same-binary A/B requires. Not a configuration |
VLLM_LTX2_DIT_FLASH_ATTN |
unset (FA-2) | THREE-WAY since #1551, because the LTX-2.5 DiT self-attention is a three-rung ladder and a two-arm knob cannot measure one. Unset selects vt::AttentionDenseFa2, the vendored FlashAttention-2 mma.sync forward, which is the default and the only rung that reaches a tensor core; =flash selects the flash-tiled scalar vt::AttentionDenseFlash, which was the #1549 default and is the denominator #1551's ratio is taken against; =0 routes back to vt::Attention, the correctness-grade vt::cuda::AttentionKernel. EACH VALUE CALLS A DIFFERENT vt:: OP, so VT_OP_PROVIDER_STATS=1 makes every arm state in its own log which rung it ran rather than leaving it inferred from the timing, which is the quantity being measured. THE UNSET DEFAULT IS THE ONLY SERVING ARM, and neither other value is a supported configuration -- retained from #1549, where this cell described a binary switch, and restated because a three-way knob has no "off": one DiT forward at 768x448/49f measures 45.547 s on GB10 through the =0 naive kernel against 6.404 s through =flash, n=119 each, one binary and one rc lease -- 7.112x (.agents/specs/ltx25-dit-attn-flash.md §10.7). That supersedes the earlier cross-run 47.84 s / ~6.0x pair, whose arms differed in binary, lease, prompt and sampler. The =0 and =flash arms also RENDER DIFFERENTLY, by more than one pixel of global image shift on three of four axes, which is sampler amplification rather than a kernel defect (#1743). The naive kernel is slow because it launches one 256-thread block per (query, head) and re-reads all K/V from global for each one -- 1.77e8 block-key iterations per call x 48 layers (issue #1549). This exists so the A/B that MEASURES the swap runs on ONE binary, which is what a same-binary A/B requires, and so a numerical regression has a same-binary control. EVERY ARM IS AN EXACT MATCH AND A FOURTH VALUE IS REFUSED BY NAME (#1751). Until then the naive arm was selected by a PREFIX (arm[0] == '0', so 0x, 07 and 0flash all ran it) while the flash arm used an exact compare, and every value matching neither -- falsh, FLASH, flash with a trailing space, naive, 1, or an empty value from an unset shell variable -- fell through to the FA-2 default in silence. A mistyped DENOMINATOR therefore ran the numerator's kernel a second time and yielded ~1.00x, which is also exactly what "no speedup" looks like, so the number could not report its own failure. Set it to anything but flash, 0 or nothing at all and the first DEVICE DiT forward now throws a message naming the variable, the value you set and the three arms. An empty value is REFUSED rather than treated as unset, because export VLLM_LTX2_DIT_FLASH_ATTN=$ARM with an unset ARM produces exactly that and is the case in which an operator most believes they selected an arm. UNSET still means FA-2: it is the shipped default and the only serving arm, so requiring the variable would refuse every production render. THE REFUSAL IS ON THE DEVICE FORWARD ONLY, because src/vllm/model_executor/models/ltx2_device.cpp is the one file in the tree that reads this variable. A --device cpu render routes to the host Ltx2DitForward (src/vllm/model_executor/models/ltx2_dit.cpp), which never reads it, so on that arm an unrecognised value is still ignored in silence. THE GAP HAS LOW CONSEQUENCE THERE, and that is why it stands rather than because nobody looked: the host self-attention call site (src/vllm/model_executor/models/ltx2.cpp) calls vt::Attention unconditionally, and on the CPU kAttention, kAttentionDenseFlash and kAttentionDenseFa2 are all registered to the same AttentionKernel (src/vt/cpu/cpu_ops.cpp), so all three rungs ARE one function and no host measurement can name a rung it did not run. What the host arm lacks is the refusal, not a correct answer, and nothing here promises to add one. A MEASUREMENT LANE in the same shape as VLLM_LTX2_POOL_DRAIN above, and never a configuration |
VLLM_RENDER_PHASE_LOG_STDERR |
unset | =1 also prints the render phase table as a fixed-width block to stderr. The table itself is written to <output_dir>/phase-log.json on the shipped default whatever this is set to (issue #1010); this only adds a copy somebody watching a terminal can read. The copy is INDEPENDENT of the file write: it is emitted whether the file was written or not, because an unwritable --output-dir is exactly the run where a console copy is the only evidence left, and a lane gated on the write would be silent in the one case it exists for. Off by default because a render that writes its table beside its frames does not need to shout, and a suite that runs ninety of them would print ninety |
VLLM_RENDER_PHASE_SAMPLER |
on | =0 stops the render phase log's 100 ms sampler thread. The phase table is still written and every phase still carries a byte count taken at its own boundaries; what is lost is the INTERIOR peak of a long phase, which is the half a minutes-long denoise needs. A MEASUREMENT LANE in the same shape as VLLM_LTX2_POOL_DRAIN above — it exists so the A/B over what the sampler costs runs on one binary — and never a configuration |
VLLM_RENDER_PROGRESS |
on | =0 silences the LTX-2.5 render's live progress lines on stderr — one [render] + <phase> when a phase opens, one [render] - <phase> ... dur= when it closes, and one [render] dit forward N phase P step k/N t=.. last=.. per DiT forward. ON by default and that is the whole point (issue #1413): the phase TABLE is written only by a generation that returns, so a render that is killed or aborted writes nothing, and between the load banner and wrote N frames a 2.5-hour render was silent — working and hung were the same observation. The open line means the LAST LINE PRINTED names the phase in flight, and last= is the per-forward cost #1375 could otherwise only get as an interval between GPU busy/idle edges from outside the process. A MEASUREMENT LANE in the same shape as VLLM_LTX2_POOL_DRAIN above — it exists so the A/B over what the emitter costs runs on one binary — and never a configuration |
VT_TTFT_DUMP |
unset | =1 prints one TTFTSPLIT rid=... intake=.. queued=.. prefill=.. decode=.. e2e=.. line per finished request to stderr, reconstructing the per-request timing split from the event-populated req_state timestamps. The async serving frontend otherwise tracks no per-request stats (passes iteration_stats=nullptr and never stamps EngineCoreOutputs.timestamp); under this flag both are wired so a serving TTFT attribution can read the queue-vs-execution split against vLLM's own request_{queue,prefill,decode}_time_seconds. Generation is byte-identical when unset (the default path is instruction-identical to production); the durable replacement is the async /metrics stat logger |
VT_LOOP_TRACE |
unset | =1 prints one LOOPTRACE ... line per ~1 s window to stderr from the engine busy loop: the full-iteration cadence (interval), process_engine_step wall (step), per-window admits, input-queue residence (resid = enqueue-to-drain, the same endpoints VT_TTFT_DUMP's intake measures), per-drain admit max and max backlog depth. Diagnoses whether the admission wait is one busy-loop iteration or the input queue is backing up (it attributed the 35B INTAKE deficit to bursty arrival during long prefill steps). Byte-identical when unset: every trace read is guarded, and the enqueue timestamp is stamped only under the flag |
VT_H3_GAUSSIAN_NOISE |
unset | minimax-h3-gen only: =1 seeds the diffusion initial noise from Box-Muller GAUSSIAN N(0,1) (what a flow-matching model is trained on) instead of the historical uniform[-1,1] draw. A/B knob for the render-coherence investigation; the exact values still do not match torch's RNG (that only selects WHICH sample), but the DISTRIBUTION does |
VT_H3_TRACE_MOTION |
unset | =1 prints one [h3-motion] step ... line per MiniMax-H3 denoise step to stderr: the step's velocity stats (v_rms/v_amax/v_mean of the DiT output), the per-step latent motion over the denoise-target rows (drows_rms), and the running latent norm (rows_rms). Because the rectified-flow Euler integration telescopes to (sigma0 - sigmaN) * v, a velocity that does not EVOLVE across steps produces a step-count-invariant result; this trace measures exactly that (added for the render-coherence bisection). Byte-identical when unset — every read is guarded and it only reads buffers the loop already holds |
VT_H3_VAE_PROBE |
unset | =1 runs a video-VAE receptive-field probe after the normal decode: it perturbs ONE interior spatial latent cell (across all channels and temporal frames), re-decodes, and prints a per-16px-block RMS-change map ([h3-vae-probe]) over output frame 0. If only the perturbed cell's block moves, the ViT3D decoder is not mixing tokens spatially. Byte-identical to production when unset (no second decode) |
VT_H3_DUMP_DIR |
unset | Directory into which the MiniMax-H3 denoise loop writes the initial and final video latent rows (init_video_rows.f32, final_video_rows.f32) and the pipeline writes the exact VAE-input latent (vae_input_video_latent.f32), all raw little-endian f32. Lets two runs (e.g. 12 vs 50 steps, conditioned vs not) be byte/stat-compared, and the video VAE decode be replayed on a KNOWN latent, without re-running the denoise. Byte-identical to production when unset (no file is opened) |
VT_H3_ACT_DUMP |
unset | =<file> writes a per-STAGE activation fingerprint of the MiniMax-H3 DiT forward (MiniMaxH3DitForwardDevice) — stats (mean/rms/absmax/finite) plus fixed positional sample values — so two weight arms running the SAME graph on the SAME inputs (e.g. the NVFP4-bf16 stream vs the FL2VA-GGUF-bf16 control) can be diffed stage-by-stage to localise where they diverge. Byte-identical to production when unset (no file opened) |
VT_H3_ACT_CALL |
0 |
Selects WHICH forward VT_H3_ACT_DUMP captures: only the forward whose 0-based call index equals this value dumps (default 0), so a single small render (--denoise-only --steps 1) captures exactly one clean forward. No effect unless VT_H3_ACT_DUMP is set |
VT_H3_DUMP_INPUTS |
unset | Directory into which the MiniMax-H3 denoise loop writes EVERY DiT input at step 0 as raw little-endian binary plus a manifest.txt — the packed layout (input_ids/image_mask/audio_mask/img_pos/audio_pos/text_pos/update_mask/cu_seqlens/document_id), the fp64 position grid (img_position_ids.f64), the per-token modality tags (token_tags.i64), the per-token pre-unique timesteps and their unique_timesteps/inverse_indices/combined_indices AdaLN selection, both sigma schedules, and the raw prompt_embeds; the minimax-h3-gen driver additionally writes prompt_token_ids.i32. Lets the REAL-scale DiT inputs be diffed EXACTLY against upstream pipeline_minimax_h3.py (the render-coherence S1 surface the reduced-dim ladder never fed real values into). Byte-identical to production when unset (no file is opened) |
Discrete-ROCm (gfx1201) bring-up knobs from PR #140. Off by default; no effect on CUDA/CPU builds beyond the documented behavior.
VT_GEMMA4_MLP_MOE_PARALLEL used to be listed here with a default and an
effect, and it was never wired: it existed only in the deferred layer-loop path
that PR #140 did not ship, so nothing in the tree has ever read the name.
Setting it does nothing, and the row is gone rather than caveated
(#2389).
| Variable | Default | What it does |
|---|---|---|
VT_GEMMA4_EXPERT_VRAM_MB |
off (unset/0) |
N>0 enables device expert LRU with N MiB fill budget; unset/0 = device expert LRU off (resident path uses VT_GEMMA4_RESIDENT_*). Evict only with VT_GEMMA4_EXPERT_EVICT=1 |
VT_GEMMA4_RESIDENT_EXPERTS |
unset | =1 preloads the Gemma-4 MoE experts resident on the GPU(s) after the first use instead of streaming them per step (discrete-ROCm optimization). No-op (with a stderr note) on a binary built without -DVLLM_CPP_HIP |
VT_GEMMA4_RESIDENT_GPUS |
2 |
Number of GPUs across which resident Gemma-4 experts are spread; clamped to the ROCm device count. Read only when VT_GEMMA4_RESIDENT_EXPERTS=1 |
VT_GEMMA4_RESIDENT_MAX_LAYERS |
(all) | Caps how many MoE layers get resident-preloaded, to fit a smaller VRAM budget. Read only when VT_GEMMA4_RESIDENT_EXPERTS=1 |
VT_GEMMA4_RESIDENT_BF16 |
unset | Force BF16 resident expert packs when set |
VT_GEMMA4_RESIDENT_NATIVE |
unset | Prefer native FP8 resident packs when set |
VT_ATTN_PREFILL_FLASH_SHAREDK |
on | ROCm Gemma-4 prefill: scoreless SharedK flash path (prod default). 0 forces DecodeGqa baseline for A/B |
VT_ATTN_PREFILL_FLASH_WMMA |
off | Lab-only ROCm flash WMMA QK path; leave off for prod |
VT_ATTN_PREFILL_SHAREDK_WMMA |
on | ROCm SharedK fused with rocWMMA QK on sliding d=256 only. d=512 stays scalar SharedK (shipping f58b WMMA spills). 0 forces scalar SharedK |
VT_GEMMA4_PREFILL_PEER_ACT |
on | Prefill MoE: run ExpertGeGLU on the expert GPU and peer activations only (not full weight PeerCopy). 0 restores weight PeerCopy |
VT_GEMMA4_GPU0_HEADROOM_GB |
12 |
GiB kept free on GPU0 when packing resident experts (decode vs long-prefill trade). Lab dual R9700 + 49k KV: 8 survives 16k+ prefill; 6 OOMs ~11k |
VT_GEMMA4_PREFILL_BATCH_MOE |
auto / 1 in lab recipe |
=1 group-by-expert prefill GEMM for T>=64; =0 serial M=1 (slow). Unset = auto |
VT_GEMMA4_DECODE_INDEXED_MAX_T |
63 |
Widen device-indexed FP8 MoE from T=1 to T<=N (N clamped [1,63]). Unset = 63. =1 restores the old T=1-only gate. T≥64 still uses prefill-batch |
VT_ATTN_PREFILL_FLASH |
off | =1 SGLang-style BM×BN GQA flash prefill (lab A/B) |
VT_GEMMA4_PREFILL_GEMM_M |
2048 |
Tokens per expert in prefill-batch GEMM chunks (16..8192; out-of-range values are ignored and the default is used). Larger M → fewer launches; lab 512 ~+37% prefill vs 64, and 512→2048 ~+80 eng @11k vs the WMMA baseline (2026-08-10), which is why the default is 2048. Lab KEEP on dual R9700 uses the default |
VT_GEMMA4_HOST_EXPERT_MB |
512 |
Host-side expert staging budget (MiB) for non-resident paths |
VT_GEMMA4_LAYER_TRACE |
off | =1 layer GPU-synced phase timers; =2 per-layer heartbeats |
VLLM_CPP_HTTP_FIXED_POOL |
1 (fixed) |
=0 reverts the HTTP worker pool to the legacy dynamic mode. Production uses the capacity-derived fixed pool; the opt-out exists for same-binary A/B attribution |
VT_ROCM_PINNED_H2D_MIB |
64 |
Chunk size, in MiB, of the ROCm bounded pinned bounce ring. A host-to-device copy of at least this many bytes, from host storage the HIP runtime does not know, to device memory, on a stream that is not capturing, is staged through four pinned buffers of this size with an event each instead of one hipMemcpyAsync off the pageable source — llama.cpp's shape at 10bf611e5, src/llama-model-loader.cpp:1440 and :1449. Host residency for the transfer is then 4 x this value for a model of any size. =0 disables the ring and restores the single call, which is the same-binary A/B. Anything that does not parse takes the default rather than 0, because a disabled ring and a working one differ only in host residency. On strix:gpu0 (gfx1151) the default is what lets the 67.56 GiB Qwen3.8-Flash-Next UD-IQ1_S produce a token at all: with =0 the same binary wedges in svm_range_set_attr and is killed at a 1200 s deadline (.agents/specs/rocm-chunked-pinned-h2d.md §7). Every other backend ignores it |
VT_ROCM_MANAGED_ALLOC |
unset | Which allocator the ROCm backend uses. Unset: hipMallocManaged only on an integrated, managed-capable device that ALSO reports PageableMemoryAccess = 1, i.e. one that can take a recoverable page fault; everything else gets plain hipMalloc. =0 never takes the managed branch. =1 takes it wherever the managed attributes allow, ignoring PageableMemoryAccess — the pre-#2511 behaviour, restored whole (allocator and unified claim), for a single-binary A/B. On an XNACK-less APU (gfx1151, gfx1103) the managed branch measured 17 GPU faults in 21 legs against 0 in 21 without it, so =1 re-arms that; a discrete card ignores all three values. Note =1 and unset also decide UnifiedMemory() there, and with it whether the CPU reference tier can serve an op ROCm has no kernel for |
VT_ROCM_ATTN_CPU_REF |
unset | =1 routes ROCm paged attention through the CPU reference kernel instead of the HIP kernel — a correctness A/B for the ROCm attention bring-up |
VT_ATTN_DECODE_GQA4 |
off | =1 routes f32-query decode through PagedAttnDecodeGqaF32Q (QG=4 fused q-heads per KV group, warp-strided walk) for bf16 or fp8-e4m3 KV at d=128/256, hq=16, kv=4; the default falls to the PagedAttnOnline reference. The arm's reduction order differs from the reference's, so greedy anchors can move at exact ties — opt-in until near-tie adjudication lands |
VT_ATTN_DECODE_SPLIT_KV |
off | =1 enables split-KV decode attention (issue #845). Uses a two-stage approach: Stage 1 computes partial attention per KV split, Stage 2 merges results via online softmax. Provides 2-8× speedup at small batch/long context by improving occupancy. Requires VT_ATTN_DECODE_KV_SPLITS, VT_ATTN_DECODE_SPLIT_WARPS for configuration. See spec .agents/specs/rocm-split-kv-attention.md |
VT_ATTN_DECODE_KV_SPLITS |
16 |
Number of KV sequence splits for split-KV decode attention (multi-CTA optimization). Each split processes a segment of the KV sequence; a reduction kernel merges partial results. Beneficial for small batch sizes and long sequences where single-CTA occupancy is low. See spec .agents/specs/rocm-split-kv-attention.md |
VT_ATTN_DECODE_SLIDE_SPLITS |
8 |
Number of KV splits for sliding-window attention decode (used when window_left >= 0) |
VT_ATTN_DECODE_SPLIT_WARPS |
12 |
Number of warps per thread block in split-KV decode stage 1 kernel |
VT_ATTN_DECODE_SLIDE_WARPS |
16 |
Number of warps per thread block in split-KV sliding-window decode kernel |
VT_DEBUG_SAMPLED |
unset | =1 prints the per-step sampled token id(s) to stderr (sampling-loop debug). Read-only; does not change output. Read once per token, so it does not stall the hot loop |
The remaining variables (roughly 120 VT_GDN_*, VT_FP4_*, VT_FP8_*,
VT_MOE_*, VT_ATTN_*, VT_NVFP4_*, VT_FUSE*, and similar) are kernel
implementation and micro-tuning switches. They are not part of the deployment
surface, their meaning is tied to a specific kernel, and they are recorded in the
engineering ledger (.agents/parity-ledger.md) and docs/BENCHMARKS.md where the
A/B that introduced each one lives. They are enumerated on
scripts/env-doc-allowlist.txt so the CI check below can tell a known
kernel-internal knob from a new, undocumented one. If you need to understand one,
grep its name under src/ for the read site and the ledger for the measurement.
scripts/check-env-doc.py runs in both directions, and a knob has to satisfy
both.
Read but not documented. It scans every VT_* / VLLM_* environment name
read from src/ and include/ and fails if any name is neither documented on
this page nor listed on the kernel-internal allowlist
(scripts/env-doc-allowlist.txt). A newly-introduced production env var therefore
cannot land silently: the author must either document it here or classify it
kernel-internal on the allowlist.
Documented but read by nothing. Every variable in a table on this page must
be read by at least one file CMake compiles (src/, include/, examples/),
and a mention in a comment does not count, so deleting the last reader of a knob
fails the gate instead of leaving its row here promising behaviour the tree does
not implement (#2389). That is
the direction the page was missing when VT_QWEN35_STAGE_MIN_FREE_FRAC and
VT_GEMMA4_MLP_MOE_PARALLEL outlived their implementations. A knob that is
legitimately unread carries a stated reason in the checker's UNREAD_EXCEPTIONS,
which is itself gated: an entry with no reason, or one whose variable has since
left the tables or gained a reader, fails the check.
The check runs in CI (the agent-record job) and has a unit/mutation test at
tests/scripts/test_check_env_doc.py.