Environment variables

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.

Deployment knobs

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
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.md
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	200000 characters	Rejects chat-completion prompts larger than this many characters. Set 0 to disable the prompt-size guard
VT_SERVER_SSE_PING_S	15	Seconds between SSE comment keepalives (:\n\n) on silent streams; <=0 disables.
VT_SERVER_MAX_NEW_TOKENS	4096	Clamps the requested generation length to this many new tokens. Set 0 to disable the cap
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_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
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
VT_SERVER_MAX_PROMPT_CHARS	200000	Rejects larger /v1/chat/completions prompts before scheduling. 0 disables the guard. This is a character count after chat-template rendering, not a token limit
VT_SERVER_MAX_NEW_TOKENS	4096	Caps the request's max_tokens value for /v1/chat/completions. 0 disables the cap

GGUF loading

Behavior of the GGUF weight loader (CPU path). See also VT_GGUF_KEEP_QUANT in the quantization format table.

Variable	Default	What it does
VT_GGUF_KEEP_QUANT	on when compute-in-quant is available	Keep GGUF weights compressed from file to matmul on CPU (no BF16 expansion), byte-identical to the reference path. 0 disables it and expands to BF16
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 in F16 rather than promoting them, an RSS/perf tradeoff
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
VT_GGUF_PREFAULT	off	Pre-fault the mmap-resident weight pages at load, trading a slower load for steadier first-token latency
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

Rollback and bisect switches

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_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
VLLM_CPP_CUDAGRAPH	on (CUDA)	Eager launches instead of a captured CUDA graph
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 $\text{sms} \times 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_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_KIMI_DEVICE_KDA_CHUNK	off (opt-in)	=1 processes the PROMPT-length KDA with the CHUNKED prefill 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 RECURRENT form, as vLLM does (kimi_gdn_linear_attn.py:141); decode stays recurrent. Spec §15 STRICT residual (c): vLLM processes the prompt chunked and we still recur, and the different reduction ORDER coin-flips the p7 near-tie. Requires VT_KIMI_DEVICE_KDA=1. Default OFF (parity-enabler; flip ON only with the token gate green)
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_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_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_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_KIMI_DEVICE_MLA	off	Kimi-Linear only: run the NoPE-MLA attention core through the shared vt::Attention device op instead of the f64 host softmax island. Requires VT_KIMI_DEVICE_COMPUTE=1. Documented negative, kept as an A/B knob: on the device-KDA config it regresses the token gate 122 to 109/128 and slows 4.24 to 3.89 tok/s, because vt::Attention's f32 online softmax is a different reduction order from vLLM's FA2 and coin-flips near-ties
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_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_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_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)
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 (only acts where the backend advertises host-addressable device memory — Vulkan today) | 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). No effect on CUDA/CPU/Metal, whose backends do not advertise the property | | 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_MAX_PROMPT_CHARS | 200000 | Hard cap on templated prompt size for /v1/chat/completions. A larger prompt is REJECTED with an error naming this variable, rather than wedging single-batch async prefill for minutes with no client tokens. 0 disables the cap | | VT_SERVER_MAX_NEW_TOKENS | 4096 | Upper bound applied to a chat request's max_tokens. Guards the same wedge from the generation side. 0 disables the cap | | 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 |

Diagnostic

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_OP_PROVIDER_STATS	off	Print per-op provider (which backend served each op) statistics
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_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
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
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)
VT_H3_ACT_DUMP	unset	Path to which ONE MiniMax-H3 device DiT forward writes a per-STAGE activation FINGERPRINT: mean/rms/absmax/finite plus a fixed set of positional sample values, for every stage and for the input-independent weight classes (islands, biases, output heads). Two weight arms running the SAME graph on the SAME inputs (e.g. an NVFP4-bf16 stream vs a GGUF-bf16 control) can then be diffed layer by layer: a JUMP at a stage names the guilty tensor class, and a scramble/transpose is caught by the positional samples even when rms matches. This is the instrument that REFUTED a #94 load-path defect (spec §8.12). Byte-identical to production when unset (no file is opened)
VT_H3_ACT_CALL	0	Which forward VT_H3_ACT_DUMP captures, as a 0-based call index within the process. Only that one forward dumps, so a single small render (--denoise-only --steps 1) yields exactly one clean fingerprint file instead of 50 overwrites. Read only when VT_H3_ACT_DUMP is set

ROCm + Gemma-4 residency (contributor #140)

Discrete-ROCm (gfx1201) bring-up knobs from PR #140. Off by default; no effect on CUDA/CPU builds beyond the documented behavior.

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 (and global-Q d=512). 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_MLP_MOE_PARALLEL	off	=1 run Gemma4 MLP and MoE on two HIP streams (lab; wall ~flat on R9700). Not wired in this PR tip (decode-graph-free split)
VT_ATTN_PREFILL_FLASH	off	=1 SGLang-style BM×BN GQA flash prefill (lab A/B)
VT_GEMMA4_PREFILL_GEMM_M	256	Tokens per expert in prefill-batch GEMM chunks (16..2048). Larger M → fewer launches; lab 512 ~+37% prefill vs 64
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_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_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

Kernel-internal knobs (deferred)

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.

Keeping this reference honest

scripts/check-env-doc.py 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. The check runs in CI (the agent-record job) and has a unit/mutation test at tests/scripts/test_check_env_doc.py.