environment-variables.md

Environment variables

This page documents the MLXCEL_* environment variables that affect mlxcel runtime, server, downloader, build, and diagnostic behavior.

Prefer CLI flags for settings that have a flag equivalent. Environment variables are useful for containers, service units, and repeatable benchmark runs, but they are process-wide and several of the low-level knobs are read once and cached on first use. Set them before starting mlxcel or mlxcel-server.

Precedence and value conventions

• If a CLI flag and an environment variable control the same option, the CLI flag wins unless the flag help states otherwise.
• LLAMA_ARG_* aliases exist for a subset of llama-server-compatible flags. This page focuses on MLXCEL_*; use --help for the full flag/env surface.
• Boolean parsing is not completely uniform across all internal knobs:
  • documented server options generally accept true/false, 1/0, yes/no, and on/off;
  • many diagnostic switches are presence-based, so any set value enables the behavior;
  • variables whose row says "falsy disables" treat 0, false, off, or no as disabled.
• Variables marked advanced or diagnostic are not a stable public API. They exist for benchmarking, rollback, or kernel-development work and may change between releases.

llama-server compatibility variables

Both server entry points accept the canonical llama-server b10621 variables below. An explicit CLI flag wins. Where a legacy mlxcel spelling is listed, the canonical variable wins over that legacy alias.

Canonical variable	Flag	Retained legacy alias
LLAMA_ARG_BATCH	--batch-size	LLAMA_ARG_BATCH_SIZE
LLAMA_ARG_UBATCH	--ubatch-size	LLAMA_ARG_UBATCH_SIZE
LLAMA_ARG_SPEC_DRAFT_MODEL	--draft-model / --model-draft / --spec-draft-model	LLAMA_ARG_MODEL_DRAFT
LLAMA_ARG_SPEC_DRAFT_N_MAX	--draft-max / --spec-draft-n-max	LLAMA_ARG_DRAFT_MAX (b10621 removed it and aborts when it is exported; mlxcel keeps it as a working fallback)
LLAMA_ARG_THINK_BUDGET	--reasoning-budget	LLAMA_ARG_REASONING_BUDGET
LLAMA_ARG_LOG_FILE	--log-file	LLAMA_LOG_FILE
LLAMA_ARG_CHAT_TEMPLATE	--chat-template	—
LLAMA_ARG_CHAT_TEMPLATE_FILE	--chat-template-file	—
LLAMA_ARG_CHAT_TEMPLATE_KWARGS	--chat-template-kwargs	—
LLAMA_ARG_THINK	--reasoning-format	—
LLAMA_ARG_REASONING	--reasoning	—
LLAMA_ARG_REASONING_EFFORT	--reasoning-effort	—
LLAMA_ARG_REASONING_PRESERVE	--reasoning-preserve / --no-reasoning-preserve	—
LLAMA_ARG_THINK_BUDGET_MESSAGE	--reasoning-budget-message (accepted, not yet injected)	—
LLAMA_ARG_SKIP_CHAT_PARSING	--skip-chat-parsing / --no-skip-chat-parsing	—
LLAMA_ARG_PREFILL_ASSISTANT	--prefill-assistant / --no-prefill-assistant	—
LLAMA_ARG_JINJA	--jinja / --no-jinja	—
LLAMA_ARG_ENDPOINT_METRICS	--metrics	—
LLAMA_ARG_ENDPOINT_PROPS	--props	—
LLAMA_ARG_ENDPOINT_SLOTS	--slots	—
LLAMA_ARG_TIMEOUT	--timeout	—
LLAMA_ARG_API_PREFIX	--api-prefix	—
LLAMA_ARG_SSE_PING_INTERVAL	--sse-ping-interval	—
LLAMA_ARG_THREADS_HTTP	--threads-http	—
LLAMA_ARG_REUSE_PORT	--reuse-port	—
LLAMA_ARG_SSL_CERT_FILE	--ssl-cert-file	—
LLAMA_ARG_SSL_KEY_FILE	--ssl-key-file	—
LLAMA_ARG_CORS_ORIGINS	--cors-origins	—
LLAMA_ARG_CORS_METHODS	--cors-methods	—
LLAMA_ARG_CORS_HEADERS	--cors-headers	—
LLAMA_ARG_CORS_CREDENTIALS	--cors-credentials / --no-cors-credentials	—
LLAMA_API_KEY	--api-key	—
LLAMA_ARG_API_KEY_FILE	--api-key-file	—
LLAMA_ARG_HF_REPO	--hf-repo	—
HF_TOKEN	--hf-token	HUGGING_FACE_HUB_TOKEN
LLAMA_ARG_OFFLINE	--offline	—
LLAMA_ARG_HF_FILE	--hf-file (always rejected)	—
LLAMA_ARG_MODEL_URL	--model-url (always rejected)	—
LLAMA_ARG_DOCKER_REPO	--docker-repo (always rejected)	—

LLAMA_ARG_TIMEOUT changed meaning in v0.7.0-beta.1 (#1432). It is now the HTTP socket read/write timeout with llama-server's 3600-second default, matching b10621. The per-request decode watchdog it used to configure moved to --decode-timeout / MLXCEL_DECODE_TIMEOUT, with its 600-second default unchanged. Setting --timeout (or LLAMA_ARG_TIMEOUT) without also setting the decode watchdog logs a migration warning at startup.

LLAMA_API_KEY and LLAMA_ARG_API_KEY_FILE are the two exceptions to "an explicit CLI flag wins" (#1437). llama-server applies every environment variable before the command line and both call the same appending handler, so the environment value ADDS a key to the set rather than being replaced by a --api-key on the command line. LLAMA_API_KEY=env-key --api-key cli-key therefore accepts both keys, on either server. Multiple keys go in one LLAMA_API_KEY as a comma-separated list, with llama-server's own splitting rules: a field may be quoted to contain a comma, and whitespace is not trimmed.

LLAMA_ARG_OFFLINE is the one variable in this table that is not bound through clap either, for a different reason (#1434). --offline carries no value, and llama-server fires a value-less option from the environment only when the value is exactly on, enabled, true, or 1; anything else, an empty value and 0 included, leaves the flag alone. mlxcel reproduces that set exactly rather than using a general boolean parser, so a variable inherited as LLAMA_ARG_OFFLINE=0 does not pin a deployment offline and a value outside the set does not abort startup.

HF_TOKEN is a credential, so --help never renders its resolved value. LLAMA_API_KEY is not bound through clap at all, so it cannot be rendered either. Everything else in this table is printed by --help the way clap normally does.

LLAMA_ARG_THINK chooses where a model's thoughts are reported (none, deepseek, deepseek-legacy, auto); LLAMA_ARG_REASONING, LLAMA_ARG_REASONING_EFFORT and LLAMA_ARG_REASONING_PRESERVE write the enable_thinking, reasoning_effort and preserve_reasoning chat-template kwargs, exactly as llama-server does, and win over LLAMA_ARG_CHAT_TEMPLATE_KWARGS when both name the same key. LLAMA_ARG_JINJA, LLAMA_ARG_REASONING_PRESERVE, LLAMA_ARG_SKIP_CHAT_PARSING and LLAMA_ARG_PREFILL_ASSISTANT are --x / --no-x pairs read at runtime, so their vocabulary is llama-server's rather than clap's. See llama-server-compat.md.

MLXCEL_REASONING_ALIAS_FIELD is the native environment equivalent of --reasoning-alias-field. It accepts reasoning (the default), which duplicates every Chat Completions reasoning_content value into an identical reasoning field, or none, which suppresses only the duplicate alias. It does not change whether a model thinks, the existing reasoning_content field, Responses API reasoning events, or Anthropic thinking blocks.

The b10621 GGML runtime options (--n-gpu-layers, --split-mode, --mlock, --numa, --rpc, the CPU thread-pool knobs, and the rest) bind their own LLAMA_ARG_* variables too (#1445). They are hidden compatibility surfaces: an inert value is accepted and anything else stops startup with a diagnostic. Value-less flags and --x / --no-x pairs among them are read at runtime rather than through clap, so their vocabulary is b10621's: a value-less option fires only on on/enabled/true/1, and a pair reads parse_bool_value plus a LLAMA_ARG_NO_* alias meaning false. See llama-server-compat.md.

LLAMA_ARG_CACHE_TYPE_K and LLAMA_ARG_CACHE_TYPE_V accept f16 from b10621's vocabulary plus mlxcel's own int8, fp16+turbo4, fp16+turbo3, turbo4 and turbo4-delegated. The other GGML quantizer names (q8_0, q4_0, q4_1, iq4_nl, q5_0, q5_1) are rejected rather than mapped onto a different quantizer, and the unquantized f32 and bf16 are rejected because mlxcel's KV cache has no f32 or bf16 storage to select.

LLAMA_ARG_CACHE_REUSE is an integer minimum reuse chunk size in llama-server, not a prompt-cache enable switch. mlxcel accepts 0 without changing prompt-cache enablement and rejects positive values with an unsupported-setting error. Use MLXCEL_PROMPT_CACHE_ENABLED to enable or disable the cache.

Google Cloud Vertex AI variables (#1456)

Both server entry points implement llama-server b10621's Vertex AI custom-container compatibility, driven purely by the AIP_* environment variables Google's platform sets (custom-container requirements). With AIP_MODE unset (or any value other than PREDICTION), nothing is registered and the other variables are ignored.

Variable	Default	Effect with AIP_MODE=PREDICTION
AIP_MODE	unset	PREDICTION (case-sensitive) enables the adapter.
AIP_HTTP_PORT	8080	Overrides --port; a warning is logged when the two differ. An unparsable value fails startup with a diagnostic.
AIP_HEALTH_ROUTE	unset	When set, mounted (leading slash ensured) as a GET alias of the health handler. Like the predict route, it is not a public endpoint: with API keys configured it requires a key, as in b10621.
AIP_PREDICT_ROUTE	/predict	The prediction route (leading slash ensured). Startup fails when it collides with a registered API route.

POST on the predict route takes {"instances": [{"@requestFormat": "chatCompletions", ...}, ...]} (at most 128 instances) and answers {"predictions": [...]} in request order. @requestFormat names the camelCase alias of any registered route (chatCompletions, completions, embeddings, rerank, messages, tokenize, ...) or a registered path verbatim; the field is stripped and the remainder is dispatched through the ordinary handler in-process, so authentication and validation apply exactly as on a direct call. A stream field is forced off with a warning, and a per-instance failure becomes an error object in that slot rather than failing the batch. See llama-server-compat.md for the manifest entries.

Common runtime variables

Variable	Values	Default	Notes
MLXCEL_DEVICE	gpu, metal, cpu	gpu hint	cpu requests CPU execution. Invalid values are ignored with a warning and treated as gpu; if no GPU backend is available, runtime falls back to CPU. The startup lines say which of the two put the runtime on the CPU: a cpu request is reported as such, next to whether a GPU backend was available, so a CPU line on a GPU host is not mistaken for a missing backend (issue #1421).
MLXCEL_WIRED_LIMIT	max, 0, none, bytes, NK/NKB, NM/NMB, NG/NGB	max	Apple Silicon GPU wired-memory limit. Unset/empty/max sets MLX's reported GPU max memory size; 0/none disables the limit; numeric values set an explicit limit.
MLXCEL_MEMORY_LIMIT	0, none, bytes, NK/NKB, NM/NMB, NG/NGB	unset	Soft MLX allocator memory cap. Unset/0/none lets MLX use its backend default; numeric values cap the allocator and make MLX raise an exception once allocations would push the working set past this value. Also feeds the mlxcel inspect / --estimate-memory preflight as the authoritative "available unified memory" figure when nonzero, through the same parse as the allocator cap (issue #1317), so 4G and 4GB produce the same preflight figure.
MLXCEL_CACHE_LIMIT	0, none, bytes, NK/NKB, NM/NMB, NG/NGB	unset	Bound on MLX's buffer cache (issue #627). Unset/0/none leaves MLX's default cache behavior; numeric values cap cached-buffer bytes via set_cache_limit. On CUDA this is the intended way to bound cache growth now that the periodic decode-loop clear is disabled by default (see MLXCEL_CACHE_CLEAR_INTERVAL): it keeps the memory pool bounded without the per-step churn that defeats CUDA-graph reuse (ml-explore/mlx#2358).
MLXCEL_CACHE_CLEAR_INTERVAL	0 (disable), positive integer (token cadence)	0 on CUDA, 256 on Metal/CPU	Cadence of the periodic clear_memory_cache in the decode loops and batch scheduler, in generated tokens (issue #627). 0 disables the periodic clear. The default is backend-aware: on CUDA the clear churns the memory pool and defeats CUDA-graph reuse (ml-explore/mlx#2358) so it is off by default (bound the cache with MLXCEL_CACHE_LIMIT instead); on Metal/CPU the cheap 256-token trim used by Python mlx-lm is kept.
MLX_CUDA_GRAPH_CACHE_SIZE (MLX-native, CUDA only)	unsigned integer capacity	2000 on CUDA builds (MLX's own default is 400)	LRU capacity for MLX's captured CUDA-graph cache. MLX keys the cache by graph shape and its lru_cache.h also counts a lifetime miss counter that never resets (not on hits, not on trim), throwing a fatal Cache thrashing runtime_error once lifetime misses pass 2 * capacity (800 at MLX's default 400). Any long-lived, shape-diverse CUDA server crosses that threshold over its lifetime, and speculative or batched decode reaches it fastest because draft/verify phases times varying batch sizes and sequence-length buckets multiply the number of distinct graph shapes; the throw is a whole-process abort, not a request-level error, so it drops every in-flight request (issue #818). mlxcel raises the default to 2000, validated sufficient (13/13 requests across bursts on GB10). This is an LRU cap, not a preallocation, so it only costs memory as distinct graph shapes accumulate. An explicit MLX_CUDA_GRAPH_CACHE_SIZE always overrides. Read only by MLX's CUDA backend, so it is a harmless no-op on Metal/CPU.
MLXCEL_HEADROOM_FACTOR	positive f64	1.20	Runtime/activation headroom multiplier used by the unified memory estimator (mlxcel inspect, --estimate-memory, --recommend-quant). Positive values <= 1.0 disable the headroom term; invalid or non-positive values warn and fall back to the default. Override only for calibration runs — see the in-code recipe in src/execution/memory_estimate.rs.
MLXCEL_CACHE_DIR	directory path	$HOME/.cache/mlxcel	Root for mlxcel's on-disk caches. The tokenizer language-analysis disk cache (language-bias features) lives under tokenizer-scripts/, and the location-independent global model store lives under models/<owner>/<name> when MLXCEL_MODELS_DIR and the store-root flag (--model-store-root on the servers, --models-dir on the subcommands) are both unset.
MLXCEL_MODELS_DIR	directory path	unset (falls back to ${MLXCEL_CACHE_DIR:-$HOME/.cache/mlxcel}/models)	Dedicated model-store root. Snapshots live directly at $MLXCEL_MODELS_DIR/<owner>/<name> with no models/ subdir, so the whole store can sit on a separate volume without dragging the tokenizer-script cache along. Read by mlxcel download, the -m/--model resolver (generate / serve / inspect / run), the mlxcel-server -m/--model resolver, and list / rm. Resolution precedence for the models root: the CLI flag (--model-store-root <PATH> on mlxcel-server / mlxcel serve since #1438 reserved --models-dir for b10621 router mode; still --models-dir <PATH> on the download / list / rm / generate subcommands), then MLXCEL_MODELS_DIR, then ${MLXCEL_CACHE_DIR:-$HOME/.cache/mlxcel}/models. (download --local-dir <PATH> is separate: it writes the snapshot verbatim at that exact path.)
MLXCEL_DEFAULT_ORG	HuggingFace org/user name	mlx-community	Org prepended to a bare, prefix-less model name (no /) by the -m/--model resolver (generate / serve / inspect / run), the mlxcel-server -m resolver, and the download verb (mlxcel download / mlx-server download), so mlxcel run Qwen3-4B-4bit resolves to mlx-community/Qwen3-4B-4bit and mlxcel download Qwen3-4B-4bit downloads that same repo. An explicit owner/name repo-id and an existing local path are unaffected. An empty/whitespace value falls back to mlx-community.
MLXCEL_SERVER_DECODE_STORAGE	auto, dense, paged	auto	Server continuous-batching decode storage. --decode-storage-backend takes precedence. Invalid values warn and fall back to auto.
MLXCEL_KV_CACHE_BUDGET	auto, unsigned integer bytes, or none/0	auto	Paged KV block-pool budget for continuous batching. --kv-cache-budget takes precedence. Defaults to auto (#628): it pairs with the batched-decode default so admission caps KV for the concurrent batch and returns backpressure instead of an OOM abort. Applies to pool-backed Fp16 caches under the paged decode backend (the --parallel > 1 default); inert on the dense backend. none / 0 leaves the pool unbounded.
MLXCEL_PREFILL_CHUNK	unsigned integer tokens, 0 disables	2048	Cache-level chunked prefill for the single-sequence mlxcel generate / mlxcel-bench-decode path: the prompt is fed through the model in chunks of this many tokens, so sliding-window KV caches trim between chunks and prefill peak memory stays near the chunk size instead of scaling with the whole prompt (issues #672/#674). 0 forces the previous single-pass prefill. Models can opt out of multi-call prefill via LanguageModel::supports_chunked_prefill. The server path is controlled by --prefill-chunk-size instead.
MLXCEL_MAX_BATCH_PREFILL_TOKENS	unsigned integer tokens, 0 disables	derived (2 * max_batch_prefill * prefill_chunk_size)	Padded-token budget bounding one server batched prefill's transient memory (issue #715). The batched path pads a cohort to its longest prompt L and materializes a [B, L, L] FP32 mask, an O(B*L^2) transient; this caps the drained window by total padded tokens (rows * L) so the mask stays within ~2*N^2 bytes, with rows past the budget spilling to the chunked single-sequence path. --max-batch-prefill-tokens takes precedence. Unset derives 2 * max_batch_prefill * prefill_chunk_size (the shipped 2 * 4 * 512 = 4096; the 2x headroom keeps a full batch of slightly-over-chunk-sized prompts in one window); 0 disables the cap (pre-#715 unbounded behavior). Only affects families with supports_batched_prefill() under --max-batch-prefill > 1.
MLXCEL_PREFILL_GRANT_INTERVAL	unsigned integer decode ticks, 0 disables	16	Prefill fairness grant for the continuous-batching scheduler (issue #1011). A prompt longer than --prefill-chunk-size admitted next to a busy decode batch runs one chunk and is then parked; this is how many consecutive decode ticks it yields before the scheduler grants it one. A C-chunk prompt therefore reaches its first token within C * (N + 1) ticks however long the batch keeps decoding, where before #1011 it waited for the batch to drain and its time to first token had no bound at all. The decoding streams pay for that bound: over one grant cycle they get N tokens per N * D + P of wall clock for a decode step D and a chunk forward P, so their mean inter-token latency during the admission window is D + P / N. Lower is faster to first token and noisier for everyone else. 0 restores the pre-#1011 arbitration exactly, unbounded wait included. An unparseable value falls back to the default rather than to 0, so a typo cannot silently reinstate the starvation. --prefill-grant-interval takes precedence. Read once at scheduler construction. Attribute a run with mlxcel_batch_prefill_grants_total on /metrics; see scripts/bench_mixed_step_admission.py and scripts/bench/starvation_probe.sh.
MLXCEL_MIXED_STEP	1/true/yes/on to enable; unset or anything else to disable	off	Experimental (issue #908). Mixed prefill/decode step prototype for the continuous-batching scheduler. Off (the default) keeps the shipped tick policy, in which a chunked prefill parked mid-prompt yields to an active decode batch until the --prefill-grant-interval grant comes due (issue #1011). On, EVERY tick decodes the active batch and then advances the parked chunk, which is the extreme end of that same frontier: the parked prompt's time to first token is as short as it can be and the decoding streams pay a chunk forward in every inter-token gap. This is the scheduling half of a mixed step, not a fused ragged forward, and it remains a measurement instrument rather than an operator knob: ADR 0005 rejects the fused execution model, and the production fairness policy shipped separately as --prefill-grant-interval. Speculative rounds still outrank mixed steps, so the two never share a tick. Read once at scheduler construction. With the variable unset, tick arbitration is the shipped #1011 policy, which differs from the pre-#908 scheduler over the complete policy state space only where the fairness grant fires. Attribute a run with mlxcel_batch_mixed_steps_total on /metrics; see scripts/bench_mixed_step_admission.py.
MLXCEL_ATTENTION_CHUNK_BUDGET_MB	unsigned integer MiB, 0 disables	1024	Advanced. Score-matrix byte budget for the CUDA attention fallback. When a configuration cannot reach a fused SDPA kernel (head_dim > 128 families such as gemma-3/gemma-4, or softcap composites) and one attention call would materialize more scores than this budget, the query axis is processed in budget-sized chunks (issue #672). Larger values trade memory for fewer kernel launches; 0 restores the unchunked fallback.
MLXCEL_ALLOWED_ORIGINS	comma-separated origin list (e.g. https://app.example.com,https://admin.example.com)	unset	Restricts CORS to the listed origins, reflecting the request Origin only when it matches one of them. --allowed-origins takes precedence, and both are mutually exclusive with --cors-origins / LLAMA_ARG_CORS_ORIGINS, which follows llama-server b10621 semantics instead (the configured string is emitted verbatim, * echoes the Origin, localhost echoes only a localhost Origin). Unset leaves the b10621 rule in force with its * default. Each value must be a bare scheme://host[:port] origin (http/https, no path or query); a malformed value fails server startup with a clear message instead of being silently dropped. Only affects the browser-reachable TCP HTTP listener: the Unix-socket transport sends no Origin header and is unaffected.
MLXCEL_ENABLE_SETTINGS_ENDPOINT	true/false, 1/0, yes/no, on/off	unset (endpoint absent)	Environment twin of --settings: opts into authenticated GET and PATCH on /v1/settings and /settings. Applied after flag parsing, so an explicit --settings on the command line wins. An unparseable value is ignored with a warning. Startup refuses to expose the endpoint on a non-loopback TCP listener without an API key; see server-features.md.
MLXCEL_ALLOW_PRIVATE_MEDIA_URLS	1 to enable	unset (refusal active)	Disables the non-public-address refusal on the server's remote media fetches (image, audio and video URLs), turning every address check into a pass, for deployments that genuinely serve their media from an internal object store. It is off by default because a fetch proxy that reaches the private network has to be an explicit operator decision. See llama-server-compat.md.
MLXCEL_ROUTER_STATS_VERBOSE	1, true, yes, on (case-insensitive)	unset (addresses redacted)	Opts the router's client-facing GET /router/stats into the full address view ("addresses_redacted": false), including each registered peer's raw host:port. By default the response redacts raw addresses and reports only the stable router-assigned node ids, roles, health status, and dispatch counts. The redaction is defense-in-depth, not a substitute for network isolation. See distributed.md.
MLXCEL_DECODE_ALLOWLIST	comma-separated numeric IP:port list	unset (permissive, logs a warning)	Opt-in defense-in-depth for disaggregated serving (issue #389): a prefill node validates the router-chosen decode_target against this allowlist before connecting for the KV-cache handoff. Set it to the FULL pool of router-selectable decode nodes; it is independent of --decode-peers, which stays the static handoff fallback only. Entries must parse as SocketAddr (hostname:port entries are skipped with a warning). Unset, or parsing to no valid entry, stays permissive and logs a warning rather than rejecting. See distributed.md and CONTINUOUS_BATCHING.md.
MLXCEL_DECODE_TIMEOUT	seconds	600	Per-request decode watchdog: cancels a request whose generation stops producing tokens. --decode-timeout takes precedence. This is the control --timeout / LLAMA_ARG_TIMEOUT carried before v0.7.0-beta.1 (#1432); that spelling is now the HTTP socket read/write timeout. 0 falls back to the built-in 300-second guard with a warning rather than expiring every request.
MLXCEL_SURGERY	YAML file path	unset	Weight-load surgery configuration path. --surgery takes precedence. Active when the surgery feature is built, which it is by default; a --no-default-features build ignores it. See Cargo feature flags.

The three size-valued variables above (MLXCEL_WIRED_LIMIT, MLXCEL_MEMORY_LIMIT, MLXCEL_CACHE_LIMIT) share one grammar, parsed in one place (parse_memory_size in src/execution/runtime.rs). Values are trimmed and case-insensitive; a K/KB, M/MB or G/GB suffix scales by the matching power of 1024 (binary units, so 1G is 1073741824 bytes and not 10^9), and no suffix means plain bytes. A suffixed value may carry a fraction (1.5GB) and the result is floored; a bare byte count stays integer-only. 0, none and an empty value mean "unset" for all three. The mlxcel inspect and --estimate-memory preflight reads MLXCEL_MEMORY_LIMIT through that same parser, so the availability figure it prints matches the cap the allocator will actually apply (issue #1317).

When mlxcel inspect --json is used, the same availability figure is exposed as raw bytes in budget_bytes; weights_bytes, kv_bytes_total, activation_bytes, headroom_bytes, and total_bytes are also raw byte counts so downstream recipe tooling can compare model and hardware fits without scraping the text banner.

Server context sizing

mlxcel serve and mlxcel-server follow llama.cpp server semantics for the llama-compatible flags --ctx-size / LLAMA_ARG_CTX_SIZE and --parallel / LLAMA_ARG_N_PARALLEL: an explicit --ctx-size C is a total context budget shared by the active request slots, so each slot receives floor(C / N) tokens when --parallel N is used. If --max-batch-size M is set, M is the divisor because it controls the maximum number of concurrent decode sequences. With --no-batch, the divisor is 1.

--parallel defaults to 4 (serving-throughput default, #628), so an explicit --ctx-size C is divided across 4 slots by default; set --parallel 1 to give a single slot the full budget. The default --ctx-size 0 (use the model's context window per slot) is not divided. Non-batching families (SSM / hybrid) run a single decode slot regardless of --parallel.

Startup fails when the effective per-slot context window is below 512 tokens. The /slots endpoint and /health.context_size report the effective per-slot window, not the total --ctx-size budget. The --estimate-memory preflight uses the same per-slot window and active-sequence count so increasing --parallel does not multiply KV memory for a fixed explicit --ctx-size.

Build-time variables

These are read by the mlxcel-core build script.

Variable	Values	Default	Notes
MLXCEL_BUILD_METAL	1/0, on/off, true/false, yes/no	on on macOS	Overrides the CMake MLX_BUILD_METAL setting for local builds. Invalid values fail the build.
MLXCEL_BUILD_ACCELERATE	1/0, on/off, true/false, yes/no	on on macOS	Overrides the CMake MLX_BUILD_ACCELERATE setting for local builds. Invalid values fail the build.
MLXCEL_CXX_MARCH	a -march= value, or none	native	ISA baseline for the C++ bridge in release builds. Set a portable baseline (e.g. x86-64-v3) for binaries that run on machines other than the build host; none omits the flag. See Installation.

CUDA builds also use non-MLXCEL_* variables such as CUDA_HOME and MLX_CUDA_ARCHITECTURES; see Installation.

OpenXLA / StableHLO backend variables

Advanced. These apply only to builds that enable the xla-backend / xla-iree Cargo features (issue #449, ADR 0004); shipping Apple-Silicon and CUDA binaries do not, so on those builds the variables below are inert and the engine is always MLX. See Installation for how to build the backend.

Variable	Values	Default	Notes
MLXCEL_BACKEND	mlx, xla	mlx	Runtime compute-backend selector read by select_backend(). xla routes load_model / forward through the OpenXLA/IREE engine; it takes effect only when the binary was built with xla-backend (else it is ignored and MLX runs). Any other value selects MLX.
MLXCEL_XLA_DEVICE	metal, cuda, local-task	metal on macOS, local-task elsewhere	IREE HAL device for the XLA engine. CUDA is never auto-selected (it needs a CUDA-enabled runtime build), so set cuda explicitly on a GB10-class host. local-task is the CPU device.
MLXCEL_XLA_VISION_BACKEND	auto, iree, host	auto	LLaVA vision-tower/projector selector for the shared CLI/server prepared-prefill path. auto tries the resident IREE module and emits the exact startup failure before using the host implementation; iree fails startup instead of falling back; host explicitly keeps vision on MLX. The selected backend and per-request IREE transfer/timing metrics are emitted through tracing.
MLXCEL_XLA_CONTEXT_CAPACITY	integer tokens in 1..=2147483647	256	Static sequence capacity compiled into the OpenXLA prefill/decode pair and every KV buffer. Requests are rejected before native execution when effective_prompt_len + max_new_tokens exceeds this value. Increasing it grows KV memory linearly and the prefill attention mask quadratically, and produces a distinct compiled artifact that may take longer to compile. StableHLO dynamic sequence shapes are not used.
MLXCEL_XLA_PRECISION	f32, f16, bf16	f32	Contraction precision the StableHLO emitter uses for matmuls (norms and softmax stay in f32). Read at graph-emit time. An explicit value forces that precision even on a GPU device whose default would differ; an unset or unrecognized value falls back to the per-device default. The committed byte-exact goldens are the f32 graphs, so a byte-exactness check rejects a non-default precision.
MLXCEL_XLA_QUANT	packed	unset	packed keeps quantized weights packed inside the graph (device-side dequant) instead of dequantizing at load. Read both at emit time and by the loader so uploaded buffers match the emitted args. A no-op on unquantized checkpoints. Not supported on the Metal target (packed int8 dequant prefill faults on the Metal HAL driver); use the CUDA or local-task target, or leave it unset to dequant at load.
MLXCEL_XLA_IREE_COMPILE	path to iree-compile	baked at build time	Runtime override for the iree-compile binary used to lower the bundled graphs to vmfbs. The CUDA source-runtime build ships no compiler, so this must point at a cuda-capable iree-compile matching the runtime version; the CPU/Vulkan dist build falls back to the dist's own bin/iree-compile.
IREE_DIST	path to an iree-dist tree	baked at build time	Non-MLXCEL_* runtime override for the IREE distribution (CPU/Vulkan build). Takes precedence over the path baked in at build time. Also the build-time variable that selects the dist to link against; see Installation.

Compiled vmfbs are cached on disk (keyed by graph text, flags, and the compiler path), so only the first load of each graph variant pays the iree-compile cost.

Downloader variables

Variable	Values	Default	Notes
MLXCEL_NO_PROGRESS	any non-empty value	unset	Suppresses interactive download progress bars. NO_COLOR and CI=true also suppress bars.
MLXCEL_ALLOW_INSECURE_ENDPOINT	any non-empty value	unset	Allows sending a Hugging Face token to a non-HTTPS HF_ENDPOINT. Leave unset outside audited internal mirrors.
MLXCEL_EXTRA_CA_CERTS	file path (PEM bundle, max 1 MiB)	unset	Adds one or more custom CA certificates to the download client's trust store (also used by the server's remote-media fetch client), on top of its default trust store. Needed behind a TLS-inspecting corporate proxy (Cloudflare Zero Trust/WARP Gateway, Netskope, Zscaler, ...) when its root CA is not available to the active TLS backend, such as in a container or scoped service environment. Point this at the proxy's root CA (concatenate multiple PEMs into one file to trust more than one). A missing, unreadable, oversized, malformed, or empty-of-certs file is a hard error.
HF_HUB_CACHE	directory path	unset	Probed read-only for an already-downloaded snapshot before mlxcel download fetches anything (the existing copy is reused, never re-fetched). Used verbatim as the HuggingFace Hub cache directory. mlxcel never writes into the HF content-addressed layout.
HF_HOME	directory path	$HOME/.cache/huggingface	Fallback HuggingFace cache root when HF_HUB_CACHE is unset; the hub lives under HF_HOME/hub. Same read-only reuse semantics as HF_HUB_CACHE.

Server prompt-cache variables

These variables are applied when the corresponding CLI flag is absent.

Variable	Values	Default	Flag equivalent
MLXCEL_PROMPT_CACHE_ENABLED	boolean	true	--prompt-cache-enabled
MLXCEL_PROMPT_CACHE_CAPACITY_BYTES	unsigned integer bytes	2147483648	--prompt-cache-capacity-bytes
MLXCEL_PROMPT_CACHE_MAX_ENTRIES	unsigned integer	1024	--prompt-cache-max-entries
MLXCEL_PROMPT_CACHE_TTL	unsigned integer seconds	3600	--prompt-cache-ttl
MLXCEL_PROMPT_CACHE_MIN_PREFIX	unsigned integer tokens	32	--prompt-cache-min-prefix
MLXCEL_PROMPT_CACHE_SNAPSHOT_CAPACITY_BYTES	unsigned integer bytes	model-aware, fallback 536870912	--prompt-cache-snapshot-capacity-bytes
MLXCEL_PROMPT_CACHE_SNAPSHOT_MAX_ENTRIES	unsigned integer	4096	--prompt-cache-snapshot-max-entries
MLXCEL_PROMPT_CACHE_SNAPSHOT_TTL	unsigned integer seconds	7200	--prompt-cache-snapshot-ttl
MLXCEL_ENABLE_VLM_PREFIX_CACHE	boolean	false	--enable-vlm-prefix-cache
MLXCEL_RESPONSES_STORE_MAX_BYTES	unsigned integer bytes	268435456	--responses-store-max-bytes
MLXCEL_CONVERSATION_STORE_MAX_BYTES	unsigned integer bytes	67108864	--conversation-store-max-bytes
APC_ENABLED	boolean	true	--apc-enabled
APC_BLOCK_SIZE	unsigned integer tokens	16	--apc-block-size
APC_NUM_BLOCKS	unsigned integer	derived from max entries	--apc-num-blocks
APC_HASH	sha256 or blake3	sha256	--apc-hash

LLAMA_ARG_CACHE_REUSE is validated independently of MLXCEL_PROMPT_CACHE_ENABLED; it is not a boolean alias for this table's enable switch.

Automatic Prefix Caching is on by default; pass --apc-enabled=false or set APC_ENABLED=false to fall back to whole-prefix matching only (a stored prefix is then reusable only when it is fully contained in the new request). The APC_* names mirror the upstream mlx-vlm env surface.

When MLXCEL_PROMPT_CACHE_SNAPSHOT_CAPACITY_BYTES and --prompt-cache-snapshot-capacity-bytes are both absent, startup may raise the 512 MiB fallback from the loaded model's config.json. The implicit default sizes six representative exact-prefix snapshots at min(context_size, 8192) tokens, including architecture-aware attention KV plus fixed recurrent state for hybrid snapshot families, and clamps that raise to one quarter of detected available memory. An explicit env/CLI value is never replaced. For Qwen3.8-27B 4-bit this avoids the old default that could fit only one roughly 500-600 MiB snapshot and then LRU-evict the live session chain.

GET /v1/cache/stats reports snapshot_bytes_per_entry and snapshot_self_evictions so operators can distinguish healthy same-session supersede from capacity thrash.

MLXCEL_ENABLE_VLM_PREFIX_CACHE opts same-image multimodal follow-up turns into prompt-prefix sharing while leaving text-only prompt-cache behavior unchanged.

MLXCEL_RESPONSES_STORE_MAX_BYTES and MLXCEL_CONVERSATION_STORE_MAX_BYTES bound the approximate retained JSON bytes for /v1/responses response history and conversation transcripts. The entry count and TTL limits still apply. A value of 0 leaves the route surface enabled but makes newly stored entries immediately evict themselves.

The three SNAPSHOT variables budget a separate store: whole recurrent-state snapshots for SSM and linear-attention families, which cannot share KV blocks and are therefore kept as exact-prefix entries. A snapshot's size tracks model width rather than prompt length, from a few MiB on a small model to 300 MB or more on a 30B-class one. The implicit model-aware default covers the common case, but explicit deployment caps should still be sized from measurement: serve one conversation, read snapshot_bytes_per_entry from /v1/cache/stats, and give the store enough headroom for the concurrent sessions and boundary/completion producers you expect. Once a turn's snapshot strictly extends the previous turn's token vector, the newer one supersedes the older within the same producer chain and snapshot_supersedes advances. If snapshot_self_evictions advances instead, capacity pressure is evicting the same session that just donated a snapshot and the capacity should be raised.

Server audio admission variables

The OpenAI audio endpoints (/v1/audio/speech, /v1/audio/transcriptions, /v1/audio/translations) dispatch work to a single dedicated worker thread over a bounded command queue. These knobs bound that queue and the per-request reply wait, so a burst of requests cannot grow memory without bound (each queued speech-to-text command holds up to the 25 MiB per-request payload) and a stuck request does not block its caller forever.

Variable	Values	Default	CLI flag	Notes
MLXCEL_AUDIO_QUEUE_DEPTH	unsigned integer	8	--audio-queue-depth	Bound on the audio worker command queue. When the queue is full, new audio requests get a structured 503 ("All slots are busy") instead of queueing without bound. A depth of 8 caps queued payload at roughly 200 MiB plus the one request in flight. A 0 is clamped to at least one queued command.
MLXCEL_AUDIO_REQUEST_TIMEOUT_SECS	unsigned integer seconds	120	--audio-request-timeout-secs	Per-request reply timeout. A stuck or pathologically slow audio request frees its blocking thread and returns a structured 504 after this, instead of hanging. The timeout does not cancel the in-flight model work on the worker; it only frees the caller. A 0 falls back to the default rather than timing out instantly.

Server embedding and reranking variables

POST /v1/embeddings and POST /v1/rerank dispatch work to separate dedicated worker threads over bounded command queues, following the same design as the audio worker. These knobs pick the side-model checkpoints, bound both queues and per-request reply waits, and size their micro-batches. See Embeddings and reranking for both endpoints.

Variable	Values	Default	CLI flag	Notes
LLAMA_ARG_EMBEDDING_MODEL, MLXCEL_EMBEDDING_MODEL	path or owner/name repo-id	unset	--embedding-model	A second checkpoint served on /v1/embeddings next to the chat model in -m; resolved like -m. The CLI flag wins over both variables, and LLAMA_ARG_EMBEDDING_MODEL wins over the MLXCEL_ alias. Combining it with an embedding checkpoint in -m is a startup error.
LLAMA_ARG_RERANKER_MODEL, MLXCEL_RERANKER_MODEL	path or owner/name repo-id	unset	--reranker-model	A checkpoint served on /v1/rerank next to the chat model in -m; resolved like -m. The CLI flag wins over both variables, and LLAMA_ARG_RERANKER_MODEL wins over the MLXCEL_ alias. Generative rerankers require this setting because they cannot be distinguished from chat checkpoints.
MLXCEL_EMBEDDING_BATCH_SIZE	unsigned integer	16	--embedding-batch-size	Texts per forward pass. Text inputs are sorted by token length and cut into micro-batches of this size, each right-padded to its longest member.
MLXCEL_EMBEDDING_MAX_LENGTH	unsigned integer	derived	--embedding-max-length	Lowers the token cap derived from sentence_bert_config.json, tokenizer_config.json and config.json (hard cap 8192).
MLXCEL_RERANK_BATCH_SIZE	unsigned integer	0 (kind default)	--rerank-batch-size	Query/document pairs per forward pass. 0 selects the reranker kind's default: 8 for text and 2 for multimodal.
MLXCEL_EMBEDDING_QUEUE_DEPTH	unsigned integer	8	--embedding-queue-depth	Bound on each embedding/reranking worker command queue; a full queue returns a structured 503. A 0 is clamped to at least one queued command.
MLXCEL_EMBEDDING_REQUEST_TIMEOUT_SECS	unsigned integer seconds	120	--embedding-request-timeout-secs	Per-request embedding/reranking reply timeout; returns a structured 504 and frees the caller without cancelling the in-flight model work. A 0 falls back to the default.
MLXCEL_EMBEDDING_POOLING	cls, mean, max, lasttoken	unset	none	Debugging override of the pooling mode resolved from 1_Pooling/config.json or the family default; logged at load. Applies to the server and to mlxcel embed.

Speculative-decoding variables

Variable	Values	Default	Notes
MLXCEL_DRAFT_KIND	dflash, mtp	auto/none	Alias for --draft-kind when the CLI flag and LLAMA_ARG_DRAFT_KIND are absent.
MLXCEL_DRAFT_BLOCK_SIZE	unsigned integer	per drafter (4 for MTP, 16 for DFlash)	Alias for --draft-block-size when the CLI flag and LLAMA_ARG_DRAFT_BLOCK_SIZE are absent.
MLXCEL_MTP_ADAPTIVE	0/false/no/off to disable, any other value (or unset) to enable	on	Adaptive B=1 MTP policy (issue #333). When on, the server profiles the first few B=1 MTP bursts of each (target, drafter, hardware, block_size) pairing (acceptance length, verify latency, drafter latency, batch size, prompt shape) and settles to a data-driven enable/decline verdict that overrides the static per-hardware gate when the measured profile is clearly favorable or unfavorable, falling back to the static default otherwise. The verdict (enable/decline plus the coarse acceptance rate, no prompt data) is persisted at ${MLXCEL_CACHE_DIR:-$HOME/.cache/mlxcel}/mtp-policy/<key-hash>.json (hint format v3; v2 hints were settled against the pre-#725 verify kernel and the pre-#736 estimator, so they are ignored and the pairing re-profiles once), so profiling runs once per pairing and a restart reuses the verdict. Changing MLXCEL_DRAFT_BLOCK_SIZE changes the block_size dimension of the key, so the old hint is discarded and profiling restarts for the new K. Set to an off value to disable profiling and use the pre-#333 static per-hardware gates. MLXCEL_ENABLE_MTP_B1 still pins the decision and, when set, suppresses profiling. The experimental batched (B>1) path is unaffected and stays behind MLXCEL_ENABLE_MTP_BATCH. The speedup estimate is measured, not modeled (issue #736): while profiling, each burst runs a couple of classic-step probe rounds (drafterless rounds whose [1, 1] verify forward is shape-identical to a classic decode step; each emits one real greedy token, so nothing is wasted and temperature-0 output stays byte-identical), and the estimator compares the measured speculative round cost (verify + drafter + walk/finalize overhead) against the measured classic step time, which is taken as the median (not the mean) of the per-burst probe means so the first burst's one-time CUDA kernel/graph compilation for the [1, 1] verify shape cannot skew the estimate toward a falsely slow classic step. This makes the verdict correct across backends and kernel eras without hardware heuristics: on GB10 with the multirow qmv verify (MLXCEL_QMV_MULTIROW, issue #725) the Gemma 4 12B pairing profiles to about 1.5× and enables, while the same pairing on the pre-#725 per-row verify profiles to about 0.5× and declines with margin. When a window collects no probe signal the pre-#736 shape heuristic (issue #638: 1.0 on Apple Silicon, sqrt(K) elsewhere) remains as the fallback. The resulting state is readable at GET /v1/internal/mtp-policy; see Adaptive MTP policy API, and read it there rather than parsing the hint files, whose format is private.
MLXCEL_ENABLE_MTP_B1	0/false/no/off to disable, any other value to force on	adaptive (per hardware)	Manual override for the singleton (B=1) MTP burst, in both directions. When set it pins the decision and disables adaptive profiling (issue #333). When unset, the adaptive policy decides (see MLXCEL_MTP_ADAPTIVE); with MLXCEL_MTP_ADAPTIVE=0 the decision is the static per-hardware default (issue #165): non-batchable targets (gemma4_unified 12B pairs, whose only decode path is B=1) default on everywhere (measured across three prompts: 1.90x to 3.14x on M5 Max, 1.74x to 2.61x on M3 Ultra, 0.95x to 1.48x on M1 Ultra, where the prose end is a 5% loss rather than a gain); batch-capable targets (the 31B + bf16 assistant) default on from Apple GPU generation 15 (M3, M4, M5), and fall back to classic decode on generation 13 (M1, M2). This was has_neural_accelerator (M5 only) until issue #1217, on the strength of ~1.2 to 1.4× on M5 Max against a ~0.75 to 0.96× regression on M1 Ultra, both measured before #1194/#1199/#1203/#1208/#1215 and neither re-measured since. M3 Ultra, which the old predicate lumped with M1 Ultra, was never measured on this pairing at all until #1217 did it: 1.95× (prose), 2.41× (source code) and 2.65× (enumeration) on 2026-08-20 under the #1215 protocol, at a verify-round cost of 1.51 classic steps against 2.71 on M1 Ultra. The discriminator is the use_qmv_wide split (an affine-quantized projection at M >= 2 runs as one wide pass from generation 15 and as K narrow passes below it), not the Neural Accelerator. M4 is grouped by that shared dispatch rather than measured. Generation 13 keeps declining and the width sweep supports that: round cost fits 0.83 + 0.170 K classic steps for this pairing on M3 Ultra against 1.14 + 0.090 K for the 12B pairing there, so the bf16 drafter costs about 1.9x as much per block position, and carrying that ratio onto generation 13's 1.35 + 0.346 K puts a block-4 round near 3.6 classic steps, which 2.96 to 3.99 emitted tokens would only just cover. See docs/benchmark_results/mtp-b1-gate-m3ultra-2026-08-20.md.
MLXCEL_ENABLE_MTP_BATCH	truthy value	off	Advanced. Forces the batched Gemma 4 MTP burst path for parity/debug testing. Not governed by the adaptive policy (issue #333), which scopes to the validated, byte-identical B=1 path.
MLXCEL_METAL4_ATTENTION	0/false/no/off to disable; unset or any other value to enable	on (where the hardware has it)	Advanced, diagnostic kill switch. Forces the M5 neural-accelerator fused attention route off on hardware that has it, so layers::metal4_causal_attention is skipped and the ordinary SDPA path runs instead. Off-switch only: it cannot turn the route on where has_neural_accelerator && macos_supports_na is false, so setting it on an M1 is inert. This route is the first suspect whenever an M5 disagrees numerically with an earlier Apple GPU generation (issue #1065), and before this switch the only way to A/B the hypothesis was to patch should_use_metal4_attention and rebuild, which is what the #1182 M5 investigation had to do. Read once per process, so set it before starting mlxcel or mlxcel-server. Inert on non-Metal builds.
MLXCEL_GDN_CHAIN_PARITY	0 to disable, any other value (or unset) to enable	on	Advanced, diagnostic escape hatch. Gates the chain-parity gated-delta Metal kernel used by Qwen 3.5 MTP's speculative verify and rollback-replay paths (issue #1165). The standard gated-delta kernel carries float32 recurrent state across a T = K verify block and rounds it to the storage dtype only once at the end, while the classic single-token decode chain rounds after every token; a T = K verify block is therefore NOT bit-identical to K consecutive single-token decode steps unless the state is rounded per in-block step. The chain-parity kernel (gated_delta_step_seqpar) does that rounding, which is what makes Qwen 3.5 MTP's temperature-0 output byte-identical to classic decode. Setting this to 0 forfeits that exactness contract, restoring the pre-#1165 block numerics for A/B attribution of the parity kernel's own cost and acceptance effect; do not set it to 0 in a deployment that needs byte-identical speculative output. Metal-only: the non-Metal ops fallback ignores the flag (the parity guarantee does not exist off Metal today). The kernel is necessary but not sufficient: byte-identity also requires every quantized projection to dispatch to the same MLX kernel at M = block_size as at M = 1, which is not true on every GPU generation or at every block width, so the runtime probe behind MLXCEL_MTP_ALLOW_INEXACT is what actually decides whether MTP engages. See docs/benchmark_results/qwen38-mtp-m1ultra-2026-08-16.md for the measured kernel cost (inside the dispatch-noise band).
MLXCEL_MTP_ALLOW_INEXACT	1/true/yes/on to enable; unset or anything else to disable	off	Engage Qwen 3.5 MTP speculative decoding even when the startup exactness probe reports that the multi-token verify block is not byte-identical to the single-token decode chain. Before enabling MTP the runtime now measures the property instead of predicting it: one synthetic verify block and the equivalent single-token chain are run from the same state on the loaded checkpoint at the configured --draft-block-size, and their logits are compared byte for byte (three independent synthetic inputs, each two short prefills plus K + 1 forwards; measured 4.9 s for the first call and 1.3 s for a later one per input on a Qwen3.8-27B 4-bit target on an M1 Ultra, the difference being MLX's one-time kernel compilation; more than one input because a kernel pair can disagree by only a byte or two out of ten thousand, at which amplitude a single draw can read as equal; memoized per (model, block width) and warmed at worker startup so it never lands on the request path). A divergence means temperature-0 speculative output would silently differ from mlxcel generate without --draft-model, so the default is to decline and run classic decode. The static conditions (Metal backend, supports_metal_gated_delta_kernel geometry) still apply and are checked first; this probe covers what they cannot, namely which MLX kernel each quantized projection dispatches to at M = K versus M = 1. That choice depends on the GPU generation, the quantization mode, the operand sizes and the block width: use_qmv_wide in mlx/backend/metal/quantized.cpp sends M >= 2 to a different reduction whenever `mode != "affine"
MLXCEL_QMV_WIDE	0/false/no/off to disable; 1 (or any other value) to pin wide	unset (wide, until the MTP gate's retry turns it off)	Operator pin for MLX's qmv_wide kernel, the faster reduction for M >= 2 quantized matmuls on Apple GPU generation 15+ (overlay in src/lib/mlx-cpp/patches/mlx/backend/metal/quantized.cpp, added by #1199). Setting the variable at all, to any value, counts as an operator pin: the MTP exactness gate's retry (retry_without_qmv_wide) is skipped in both directions, so MLXCEL_QMV_WIDE=1 keeps the wide kernel and makes a failing probe decline MTP instead of buying exactness back, and MLXCEL_QMV_WIDE=0 runs the whole process narrow from the start. Unset, the kernel is wide until a failing MTP probe's retry finds the narrow kernel exact and pins the process narrow for good. The pin is process-wide and sits on the dispatch path of every quantized matmul; what non-MTP work pays for the narrow state is measured in benchmark_results/qmv-wide-pin-tax-m3ultra-2026-08-22.md (nothing measurable on batched decode, about 15 ms per prompt-cache-hit request's suffix prefill). Read once per process at first dispatch; mlxcel_core::set_qmv_wide can move it at runtime and the gate is its only caller.
MLXCEL_MTP_BLOCK_CONTROLLER	proxy to pin the acceptance-proxy gate; requested to honour the requested width immediately (measurement mode); any other value (or unset) selects the throughput comparator	throughput	Which controller decides the B=1 MTP verify width when --draft-block-size exceeds the drafter's configured depth (issue #1207). The default measures the decision: after a short warm-up the round loop alternates measurement windows (32 rounds) between the configured depth and the requested ceiling, compares emitted tokens per millisecond of round time, holds whichever measures faster (a challenger needs a 2% lead; ties go to the configured depth), and re-challenges the loser on a growing backoff (4, 16, then every 64 windows), with a collapsing challenger aborted after 4 rounds once it trails by more than 35% so a harmful ceiling (the Qwen 3.8 pairing measures 5.80 against 21.30 tok/s at width 12) costs rounds rather than windows. Evidence lives in the generator, so a server process keeps it across requests; drafters that set prefer_requested_block_size (Qwen 3.5 MTP) bypass both controllers and always honour the request, exactly as before. Set to proxy to restore upstream's fully-accepted-prefix gate, which issue #1207 measured holding the Gemma 4 12B pairing about 5% below its optimum (93.54 against 98.16 tok/s at requested width 5 on an M5 Max) because at 0.585 acceptance the configured prefix is rarely fully accepted no matter how profitable widening is. The batched (B>1) loop stays on the proxy gate regardless: the row-averaged accept length is the only per-round signal it measures today. Set to requested to bypass both controllers and draft at the requested width from the first round: this is the width-sweep measurement mode (#1207's own sweep needed a temporary code patch to hold widths; this value is that patch as a switch), not a deployment setting, and the harmful-ceiling protection is off under it. Read once per generator.
MLXCEL_MTP_TICK_SLICE	0/false/no/off to disable, any other value (or unset) to enable	on	Tick-cooperative B=1 MTP serving (issue #734). When on (the default), a B=1 MTP request on the Gemma 4 family is served one speculative round per scheduler tick, alternating with the classic decode/prefill actions, so concurrent classic-decode rows advance between rounds and the head-of-line stall a speculative request imposes drops from the whole burst to about one round (burst_wall_ms in the finalize log reports the max single-tick wall). Tokens stream per round instead of in one end-of-burst lump. Set to an off value to restore the legacy run-to-completion burst (the whole request served inside one tick). The interleaving trades roughly 27% of the speculative request's own aggregate decode throughput (cross-tick round gaps) for that bounded stall, so a deployment serving speculative requests without concurrent classic traffic can turn it off to keep the full-throughput burst. Greedy output, acceptance accounting, and every other env gate are unchanged in both modes; DFlash and the batched B>1 paths always run to completion regardless of this flag.
MLXCEL_MTP_SLICE_GRANT_ROUNDS	non-negative integer	8	Grant budget for one hold of the tick-slice speculative slot (issue #746), counted in executed slices (slice 0, the prefill + seed, counts as the first slice of a grant). While a slice is in flight, up to 2 further tick-slice-eligible requests park in a grant backlog instead of permanently falling back to classic decode; once the active request has run this many slices with the backlog non-empty, it parks at the next round boundary and the slot is granted to the next request (priority lane first, FIFO within a lane, with an anti-starvation floor: an entry passed over by 2 grant decisions is granted next regardless of lane), so concurrent long streams share speculative acceleration in bounded turns. The budget is read once per grant and per admission decision (cached for the per-round expiry check), so changes apply from the next grant. The budget binds only under contention: a single speculative request never rotates and behaves exactly as under #734. Rotation preserves per-request token streams byte-identically (the drafter is re-armed from the session's own stored verify output at every round). 0 disables rotation and restores the pre-#746 behavior: the active request holds the slot for its whole generation and every concurrent speculative request falls back to classic decode. Unparseable values fall back to the default.
MLXCEL_SPECULATIVE_STOCHASTIC_ACCEPT	1/true/yes/on to enable; unset or anything else to disable	off	Acceptance-optimal speculative acceptance for the classic SpeculativeGenerator path (offline mlxcel generate --draft-model), issue #902. When on, temperature > 0 verification uses modified rejection sampling (accept the drafted token t iff u * q(t) <= p(t) for a fresh u ~ U[0,1), and on the first rejection emit a draw from the normalized residual relu(p - q)) instead of the default sampler-match rule (accept iff the draft equals an independent draw from the target sampler). Both rules are distribution-preserving: the emitted stream is a target-only sample either way, which is the central correction to the issue's premise. What changes is the acceptance probability, which rises from sum_x p(x) q(x) to sum_x min(p(x), q(x)), the maximal-coupling ceiling for any correct rule. Opt-in rather than default because the gain is the ratio between those two quantities and it collapses toward 1 whenever the drafter is confident (q(t*) ~ 1 makes min(p, q) and p * q coincide); measured at about 1.02 on a Llama-3.1-8B / Llama-3.2-1B pair at temperature 0.7, which does not pay for two extra full-vocabulary passes and a host sync per verified position. Check the available gain with MLXCEL_SPECULATIVE_ACCEPT_DIAG=1 before enabling. Enabling changes the RNG stream, so at an equal seed the emitted tokens differ from a default run even though the distribution is identical. Greedy (temperature == 0 or top_k == 1) never reaches either rule and is byte-identical. The Gemma 4 MTP and DFlash round loops are unaffected: they select the target token by argmax regardless of temperature, so this switch is inert there. SpeculativeGenerator::with_stochastic_acceptance(bool) overrides it programmatically. Read once per process. See speculative-acceptance.md.
MLXCEL_SPECULATIVE_ACCEPT_DIAG	1/true/on/yes to enable; unset or falsy to disable	off	Per-position acceptance diagnostic for the classic speculative path (issue #902). Adds a second stdout line after a speculative run reporting closed_form_sum_min (sum_x min(p(x), q(x)), the probability modified rejection sampling accepts) and closed_form_sum_prod (sum_x p(x) q(x), the probability the pre-#902 rule accepts), both averaged over the verify positions the run actually tested. Because min(a, b) >= a * b for a, b in [0, 1], sum_min >= sum_prod holds for every pair of distributions, so the two numbers turn an acceptance-rate anomaly into an arithmetic statement: the measured per_position_acceptance sits at sum_prod under the default sampler-match rule and at sum_min under the opt-in acceptance-optimal rule, and their ratio is the upper bound on any throughput gain the change can produce for that model pair. Enabling it also makes the argmax arm capture the drafter's proposal distribution so both halves of an A/B report the same quantity. Costs two full-vocabulary reductions and a host readback per tested position, which is why it is off by default. Read once per process. See speculative-acceptance.md.

Generation loop detection (issue #432)

N-gram tail repetition detection ends a generation early when the raw generated token stream collapses into a short repeated pattern (a single token such as 様様様様, or a short block such as abcdabcd...). It runs on the raw stream, so it also catches loops inside the reasoning/thought channel and tool-call JSON, not just the final answer. The wire finish_reason is stop, the same as vLLM. Sampling penalties (repeat_penalty, DRY) cannot recover once the logits collapse, which is why this is a stop condition rather than a logit reshaper.

The detector mirrors vLLM's SamplingParams fields, with the same JSON names on the OpenAI chat surface:

Field	Meaning
max_pattern_size	Largest N-gram pattern size to scan. 0 (default) disables detection.
min_pattern_size	Smallest N-gram pattern size to scan. 0 (default) is treated as 1; clamped to <= max_pattern_size.
min_count	Minimum consecutive repeats of a pattern that ends generation. Must be >= 2; any smaller value disables detection.

The preferred activation surface is engine-level: detection is default-on for the Gemma 4 family with no configuration required on tool-shaped requests, so a downstream serving app needs no setup for the protected traffic. The gate is satisfied by tool declarations that reach the rendered prompt (tools non-empty and tool_choice other than none) or by tool-shaped message content (any message carrying tool_calls or tool_call_id, as an agent loop replaying prior turns sends). Plain chat, plain completion, and grammar-only structured-output requests keep the disabled baseline. Grammar constraints are excluded because a schema can legitimately require or permit a long run of identical values, which token-level detection cannot distinguish from a collapse. A collapse without the tool-shaped signal is still possible, since the defect is upstream in the weights; MLXCEL_LOOP_DETECTION=on and the per-request fields below remain available and are not gated.

Variable	Values	Default	Notes
MLXCEL_LOOP_DETECTION	off/0/none/false/disabled, on/default/true/enabled, or MIN,MAX,COUNT (also MIN:MAX:COUNT)	unset	Global operator override for any model. Unset lets the Gemma 4 family default-on apply to tool-shaped requests (everything else stays disabled). off force-disables for every request, including the Gemma 4 family; on forces the recommended threshold 1,20,12 for every model and every request, plain chat and structured output included; an explicit triple (e.g. 1,20,12) sets exact values. Set 1,20,4 to restore the pre-#967 threshold. A malformed value warns and is ignored.

Resolution precedence, highest first:

1. Explicit per-request fields. If a chat request sets any of max_pattern_size / min_pattern_size / min_count, those values are used verbatim, including an explicit disable (max_pattern_size=0). A client never has to send anything; the fields are only for tuning or opting out.
2. Global override. MLXCEL_LOOP_DETECTION, which an operator can use to force-enable, tune, or force-disable for any model.
3. Gemma 4 family default-on, gated on the request. When the loaded model is in the Gemma 4 family (Gemma4, Gemma4VLM, Gemma4Unified) and the request is tool-shaped, the conservative threshold min_pattern_size=1, max_pattern_size=20, min_count=12 is applied. Tool-shaped means either declarations that reach the rendered prompt or messages carrying tool_calls / tool_call_id. The declarations half is read from the same helper the chat template uses to pick which tools to render, so tool_choice: "none" alone does not count: it renders no declarations, leaving a prompt identical to plain chat. It does not disarm a turn that replays tool calls, though, since those reach the prompt anyway.
4. Disabled. The default for every non-Gemma-4 model and for every request without a tool-shaped prompt, including grammar-only structured output, preserving the exact baseline output.

KV cache and TurboQuant variables

Use CLI flags such as --cache-type-k, --cache-type-v, --kv-cache-mode, --turbo-boundary-v, and the batch KV quantization flags when possible. The variables below are useful for service-level defaults and A/B experiments. See TurboQuant KV cache for the user-facing mode descriptions. For model-state snapshot families, non-FP16 attention KV modes are not silently serialized into snapshot entries today: donation logs a named warning and skips snapshot reuse until the family has sidecar snapshot support. The requested mode remains visible through startup logs and kv_cache_mode_effective.

Variable	Values	Default	Notes
MLXCEL_KV_BOUNDARY_V_LAYERS	integer count	2	Number of first/last layers kept at higher precision for Turbo4-family modes. 0 disables. --turbo-boundary-v writes this value before cache construction and takes precedence.
MLXCEL_TURBO_BOUNDARY_V	integer count	fallback alias	Compatibility alias for MLXCEL_KV_BOUNDARY_V_LAYERS; the primary name wins when both are set.
MLXCEL_KV_SKIP_LAST_LAYER	boolean	true	Fallback for --kv-skip-last-layer in continuous-batching KV quantization.
MLXCEL_SPARSE_V_THRESHOLD	non-negative float	1e-6	Sparse-V alive threshold. 0 disables sparse-V; invalid values warn and use the default.
MLXCEL_SPARSE_V_KERNEL	falsy disables	enabled on macOS	Allows the fused Sparse-V/dequant Metal kernels. Set 0, false, off, or no to force graph fallback.
MLXCEL_SPARSE_V_COUNT	output file path	unset (off)	Diagnostic. When set to a non-empty path, every single-token Turbo4Asym decode appends call_idx,kv_tokens,skipped,total to that CSV, counting post-softmax attention weights below MLXCEL_SPARSE_V_THRESHOLD across all heads. Aggregate per layer offline by grouping on kv_tokens (constant within a decode step) and using each call's position within the step as the layer index. A graph-only side computation; zero cost when unset. See scripts/measure_sparse_v_skip_rate.sh.
MLXCEL_TURBO4_DEQUANT_SDPA	falsy disables	on	Controls the dequant-first SDPA path for symmetric Turbo4.
MLXCEL_TURBO4_ASYM_DEQUANT_SDPA	falsy disables	on	Controls the dequant-first SDPA path for asymmetric Turbo4Asym (FP16 K + 4-bit V). Falsy values fall back to the lossy sparse-V approximation.
MLXCEL_TURBO4_DELEGATED_DEQUANT_SDPA	falsy disables	on	Controls the default dequant-first SDPA path for Turbo4Delegated.
MLXCEL_TURBO4_DELEGATED_FUSED	truthy enables	off	Advanced. Enables the older custom fused delegated-kernel route, mainly for comparison when dequant-first SDPA is disabled.
MLXCEL_TURBO4_DELEGATED_FP16_FAST_PATH	truthy enables	off	Advanced. Keeps a unified FP16 V working set in delegated mode for speed experiments while maintaining packed sidecars.
MLXCEL_TURBO4_DELEGATED_FP16_SIDECARS	predecode, eager, lazy, on-demand	predecode	Sidecar maintenance policy for the delegated FP16 fast path.
MLXCEL_ENABLE_DIRECT_PREFILL_CACHE_STORE	presence enables	off	Advanced. Installs the incoming prefill tensor directly as the initial KV cache buffer when applicable.
MLXCEL_MLA_ABSORBED	1/true/on/yes enables	off	DeepSeek-family matrix-absorbed MLA decode. Caches the compressed latent (ckv, kpe) instead of the decompressed per-head K/V, cutting the KV cache from num_heads * (qk_head_dim + v_head_dim) to kv_lora_rank + qk_rope_head_dim bytes per token per layer. Costs fixed weight memory: kv_b_proj is dequantized at load and kept dense. Currently wired for deepseek_v2; deepseek_v3 / deepseek_v32 already absorb unconditionally and ignore this. Declines (and keeps the decompressed path) for any non-FP16 KV cache mode and for paged-backed caches. Prints one stdout line at load stating how many layers folded. See MLA absorbed decode.
MLXCEL_MLA_SPLIT_KV	1/true/on/yes enables	off	Advanced. Cuts the latent range into chunks whose partial softmax states are merged by the issue #898 merge kernel. Requires MLXCEL_MLA_ABSORBED; ignored without it. The partial producer is currently composed from MLX ops, so this is a correctness path rather than a speed path.
MLXCEL_CASCADE_ATTENTION	1/true/on/yes enable; 0/false/off/no kill switch	off	Two-level cascade decode (issue #903): a whole-page prompt prefix shared by several concurrent sequences is attended once for the subgroup instead of once per sequence, and merged into each member's per-request suffix state with the issue #898 merge kernel. Detected from the paged page table, where two requests naming the same physical row at the same position hold one refcounted block, which is what an APC prefix adoption produces. No kernel is added: both levels are ordinary paged-decode v2 launches. Off by default because the throughput benefit has not been measured on a serving workload yet; when unset the decode path does no page-table scanning and the flat launch is byte-identical to pre-#903. Sequences outside the subgroup, sliding windows trimmed into the middle of a page, soft-cap families and multi-token steps all keep the flat path. See Cascade attention.
MLXCEL_CASCADE_MIN_SHARED_PAGES	non-negative integer	16	Whole pages a subgroup must share before cascade is used; 16 pages is 512 tokens at the default block size of 32. Below this the two extra launches and the merge cost more than the duplicated reads they remove. 0 disables cascade outright. A value that is not a non-negative integer is ignored in favour of the default.
MLXCEL_CASCADE_MIN_MEMBERS	integer >= 2	2	Sequences that must share the span. Below 2 there is no duplication to remove and cascade is declined.

Paged decode v2 variables

Fused paged-attention decode v2 (issue #899) is the production decode path for pool-backed continuous batching, so these are operator-facing knobs rather than kernel-development switches. None of them is normally needed. Use them to size the pool for a longer context than --ctx-size implies, or to re-measure the dispatch floors on new hardware.

The kill switch is MLXCEL_PAGED_ATTENTION_NATIVE (documented under Hardware and kernel diagnostic variables): a force-off value restores the pre-#899 gather path end to end, and a force-on value pins v2 for every servable shape, bypassing the floors below. For the dispatch policy, the per-outcome startup log lines, and what the fused kernel declines to serve, see Continuous batching.

Variable	Values	Default	Notes
MLXCEL_PAGED_SLAB_BLOCKS	positive integer blocks, used verbatim (neither floored nor capped); 0 pins the pool default	derived (ceil(per_slot_ctx / block_size) * batch, floored at the 32-block pool default and capped at the per-layer share of the KV block budget)	Paged pool slab size in blocks. The fused kernels read one contiguous buffer per side, so a layer whose rows outgrow its first slab silently returns to the gather path; the server derives the slab from --ctx-size and --parallel and logs the resolved value at startup (Paged KV slab size: N blocks per layer). Raise it when serving contexts longer than --ctx-size implies. 0 pins the historical 32-block pool default, which keeps pre-#899 allocation behaviour and, as a side effect, keeps the fused path unreachable for anything past one slab. A value that is not a non-negative integer warns and is ignored, so the derived size applies exactly as if the variable were unset (#1137).
MLXCEL_PAGED_V2_MIN_KV_TOKENS	non-negative integer tokens	4096	Visible KV tokens a lone-request launch must reach before v2 is dispatched. 4096 is the weakest measured single-request win (1.08x); batch 1 at 1024 tokens is the only measured loss (0.91x), which is what the floor excludes. An unparseable value falls back to the default, so a typo cannot silently install an unmeasured policy.
MLXCEL_PAGED_V2_MIN_KV_TOKENS_PER_REQUEST	non-negative integer tokens	512	Visible KV tokens per request a multi-request launch must reach. The smallest measured multi-request win is batch 4 at 1024 tokens per request (1.41x); the floor sits below that with margin so a prompt landing just under a nominal 1K is not declined for no measured reason. Same fallback-on-typo behaviour as the row above. To benchmark the declined regime, force the path on with MLXCEL_PAGED_ATTENTION_NATIVE=1 rather than driving a floor to zero.

Video and local-media variables

These apply to video-capable VLM request handling.

Frame extraction shells out to the system ffmpeg and ffprobe, which must both be on PATH and must be ffmpeg 5.0 (2022) or newer. See supported-models.md for what that floor is and why.

Variable	Values	Default	Notes
MLXCEL_VIDEO_DIR_ALLOWLIST	comma-separated directories	unset	Local video_url file paths are rejected unless they resolve under one of these canonicalized directories. Keep directories owner-writable only; group/world-writable entries warn at startup.
MLXCEL_VIDEO_MAX_PIXELS	unsigned integer	16777216	Rejects source videos whose width × height exceeds the cap.
MLXCEL_VIDEO_MAX_DURATION_SEC	float seconds	600	Rejects source videos longer than the cap.
MLXCEL_VIDEO_MAX_PNG_FRAME_BYTES	unsigned integer bytes	268435456	Per-frame cap for the ffmpeg PNG stream splitter.

Hardware and kernel diagnostic variables

These variables are for profiling, rollback, or experiments. They are not recommended as normal deployment settings.

Variable	Values	Default	Purpose
MLXCEL_NO_PADDED_PREFILL	presence disables	auto	Disables M5+/Neural-Accelerator prefill tile alignment.
MLXCEL_FORCE_PADDED_PREFILL_MASK	presence enables	off	Forces an explicit padded prefill mask path for debugging.
MLXCEL_LOG_NA_ATTENTION	sampled, all, truthy	off	Logs Neural Accelerator attention dispatch decisions.
MLXCEL_ENABLE_FUSED_CAUSAL_PREFILL_ATTENTION	presence enables	off	Enables an experimental Llama-family fused causal prefill path when supported. Ignored for a checkpoint whose config.json sets rope_traditional, or whose rope_scaling block selects a frequency table or a position scale: the launcher applies the split-half rotation from a plain rope_theta inside C++ and takes no flag, so such a model always uses the graph path and a one-time notice on stderr says so.
MLXCEL_ENABLE_FUSED_QKV_SPLIT_ROPE	presence enables	off	Enables an experimental fused QKV projection/split/RoPE path. Ignored for a rope_traditional or rope_scaling checkpoint for the same reason as the row above.
MLXCEL_GEMMA4_ENABLE_FUSED_QKV	presence enables	off	Enables a Gemma 4 fused-QKV projection experiment.
MLXCEL_BAILING_LINEAR_CHUNKED_PREFILL	0/false/off/no disable; any other value or unset enables	on	Chunked closed-form evaluation of the bailing_moe_linear (Ling / Ring) GLA recurrence at prefill, over 64-token chunks. Default-on as of #1040: measured 2.1x to 2.4x faster prefill on Ring-mini-linear-2.0-4bit (median of 5 on an M1 Ultra: 3.85s to 1.61s at ~512 prompt tokens, 5.79s to 2.81s at ~2048, 18.17s to 8.54s at ~8192) and lower teacher-forced perplexity at every window length measured (115.83 to 114.09 at 128 tokens, 39.39 to 36.39 at 1024, 85.96 to 74.14 over 16384 in 512-token windows), because the intra-chunk sum lands in a matmul accumulator instead of a bf16 running state that compounds its error over the sequence. Decode is unchanged by construction, and confirmed: 142.9 against 143.6 tok/s. The two paths decode different continuations from one checkpoint, so set this to 0 for upstream's sequential recurrence when diffing against mlx-lm.
MLXCEL_V4_FLAT_INDEX	presence forces the flat scan	off (hierarchy when the size gate opens)	DeepSeek-V4 diagnostic (issue #549). Forces the sparse layers' pooled-row selection onto the flat O(Np) scan, disabling the two-stage HiSA hierarchy while leaving sparse attention itself engaged. Both paths score the same pooled rows the same way, so running one prompt each way separates "the hierarchy retained the wrong blocks" from "the sparse combine is wrong", which is otherwise hard to bisect because the hierarchy only engages past index_block * index_keep (1024) pooled rows, roughly 4096 prompt tokens at ratio 4. Costs memory: the flat scan materialises the [B, H, L, Np] score tensor the hierarchy exists to avoid.
MLXCEL_V4_DENSE	presence forces dense pooled attention	off (sparse past index_topk)	DeepSeek-V4 diagnostic (issue #549). Keeps every sparse layer on the dense local+pooled concat past index_topk instead of gathering the selected rows, making the layer the full-attention oracle the sparse path approximates. Below index_topk this is already what runs, so it only changes long-context behaviour. The indexer's compressor still runs and advances its pooled cache, so removing the variable resumes sparse selection with a correctly populated pool. Not a serving mode: the concat carries all Np pooled rows, so memory grows with context exactly as the sparse path was built to avoid.
MLXCEL_DISABLE_COMPILED_SWITCH_QGEGLU	presence disables	compiled path on when supported	Rolls back Gemma 4 compiled Switch-QGeGLU decode path.
MLXCEL_NVFP4_DENSE_REPACK	1/true/on/yes (matched case-insensitively) forces the dense fallback; unset or any other value keeps the default	off (direct transcode)	Any build (route only). Forces the older dense f16 repack route instead of the direct ModelOpt-triplet transcode (issues #693/#705) when loading ModelOpt NVFP4 checkpoints. On CUDA, and on non-CUDA when MLXCEL_NVFP4_NATIVE_REPACK=1 is also set, the dense route targets MLX native NVFP4; on plain non-CUDA it targets the affine 4-bit fallback, preserving the pre-#705 comparison/rollback path. Debug/parity fallback only: the direct transcode is bit-exact to the checkpoint, while the dense repack re-derives block scales and drifts by roughly one FP8/FP4 rounding step. Wins over the direct route when set.
MLXCEL_NVFP4_NATIVE_REPACK	1/true/on/yes (matched case-insensitively) selects native NVFP4 inside the dense fallback; direct native is already the default	default native on every build	Compatibility/rollback selector (issues #694/#705). Direct ModelOpt-triplet transcode is now the default on Metal/CPU as well as CUDA. This variable is retained as a compatibility no-op for the direct route and as the way to steer MLXCEL_NVFP4_DENSE_REPACK=1 toward dense f16 -> native NVFP4 instead of the plain non-CUDA dense-affine rollback.
MLXCEL_ENABLE_SOFTCAP_GQA_DECODE_GROUPED	any value except 0 enables, 0 disables	on	Grouped softcap-GQA decode. At q_len == 1 with no mask it keeps K and V at [B, H_kv, S, D] and broadcasts n_rep inside the matmul; the fallback expands them first, writing an n_rep-sized copy of the whole live cache on every decode step. On by default since #1686 measured gemma2-2b-4bit decode 195.59 to 221.69 tok/s and gemma-2-9b-8bit 40.51 to 47.34 on M5 Max at 512 prompt tokens, with greedy output byte-identical either way. Only softcap families reach this path at all (Gemma 2 sets attn_logit_softcapping).
MLXCEL_DISABLE_SOFTCAP_GQA_DECODE_GROUPED	1 disables, 0 enables	unset	Legacy rollback knob for grouped softcap-GQA decode, checked before the enable variable above. Set to 1 to restore the repeat-based path.
MLXCEL_DISABLE_SINGLE_QUERY_MASKLESS	truthy disables	maskless path on	Disables the single-query maskless attention path.
MLXCEL_EXPERIMENTAL_BOOL_CAUSAL_MASK	truthy enables	off	Enables an experimental boolean causal-mask path.
MLXCEL_PAGED_ATTENTION_NATIVE	1/true/on/yes force the native kernel; 0/false/off/no force gather (case-insensitive); unset or any other value defers to the adaptive selector	selector-governed (server: the issue #899 token floors; library entry point on Metal: native only for batch>=4 + ctx<=4096 + single-slab; on CUDA: native for any single-slab layer since #634)	Overrides the fused paged-attention decode kernel (Metal since epic #116 Phase 6/#123, and since #634 also CUDA via mx.fast.cuda_kernel). Two consumers, per resolve_dispatch_decision and resolve_paged_v2_dispatch in src/lib/mlxcel-core/src/layers.rs: (1) the library-only paged_decode_attention_pooled entry point, where since #331 an unset value no longer means "always gather" (select_pooled_paged_dispatch picks the kernel only inside the regime ADR 0001 measured it winning) and this variable force-pins either arm for A/B testing; and (2) the server's pool-backed batched paged decode, where issue #899 made the fused v2 kernel the production path and named this variable's force-off values its kill switch: setting one restores the pre-#899 gather-then-SDPA behaviour end to end, and a force-on value pins v2 for every servable shape, bypassing the measured token floors (MLXCEL_PAGED_V2_MIN_KV_TOKENS, MLXCEL_PAGED_V2_MIN_KV_TOKENS_PER_REQUEST). The override is checked before the selector in both consumers, and a forced dispatch still goes through the kernel's structural declines (single-slab layer, servable geometry, non-empty batch). History: #710 retired the library entry point to a library-only API, which is where the "not a server knob" reading came from; #899 gave the variable its second, server-side consumer. The separate DecodeBatchContext::use_native_paged_kernel request still governs the dense-compat block-table decode, a different path. See ADR 0001 and Continuous batching.
MLXCEL_PAGED_ATTENTION_V2	1/true/on/yes enable (case-insensitive, trimmed); unset or anything else disables	off	Diagnostic. Issue #898's comparison gate: makes the library-only v1 entry point (PagedBlockPool::paged_decode_fused, reached from paged_decode_attention_pooled) try the fused v2 kernel first. It does not gate the server's production v2 decode, which is issue #899, is default-on, and is governed by the token floors and MLXCEL_PAGED_ATTENTION_NATIVE instead. Read once per process, and a typo degrades to v1 rather than enabling a path nobody asked for. See ADR 0001.
MLXCEL_PAGED_DECODE_V2_CHUNK	positive integer pages	unset (planned or autotuned chunk size)	Diagnostic. Pins pages_per_chunk for every paged decode v2 launch, overriding both the plan's occupancy heuristic and any autotuned value (issue #906), matching the MLXCEL_PAGED_DECODE_SPLITS / MLXCEL_QMM_TILE_* escape-hatch convention. The plan clamps the value into the feasible range, so an over-large value is safe; a non-positive or unparseable value warns and is ignored in favour of the planned size. Read once per process.
MLXCEL_PAGED_DECODE_V2_TARGET_CTAS	positive integer	derived (Apple: gpu_core_count * 8, floored at 64; every other host including CUDA: 512)	Diagnostic. Moves the occupancy target the v2 plan binary-searches its chunk size against, as opposed to MLXCEL_PAGED_DECODE_V2_CHUNK, which pins the chunk size itself. This is the knob to turn when the derived device parallelism is wrong. The Apple figure is the gpu_core_count device-scale proxy, not a calibrated core count, and the fixed non-Apple target is unvalidated on CUDA (a CUDA-specific target should come from the SM count), so both are starting points. A non-positive or unparseable value warns and uses the derived target. Read once per process.
MLXCEL_SPARSE_PAGED_ATTENTION	0/false/off/no disable; any other value or unset enables	on	Kill switch for fused sparse decode via page indirection (#904). When disabled, a sparse-attention family falls back to the gather or additive-mask path it used before. Read once per process; an unrecognised value enables, so a typo leaves the measured path in place. Currently reached by MiniMax-M3 block-sparse decode; see sparse-paged-decode.md.
MLXCEL_SPARSE_PAGED_DUMP	1/true/on/yes enable	off	Prints each request's selected pool rows for a fused sparse decode (#904). Synchronizes with the device, so it is a debugging aid only and must never be set for a timed run.
MLXCEL_SPARSE_PAGED_MIN_SPARSITY	non-negative integer	8	Minimum live_len / selected_rows ratio a launch must clear before fused sparse decode is dispatched (#904). The fused kernel measured 0.67x at 2x sparsity and 1.22x at 8x against tuned dense SDPA, so a token floor alone is not sufficient. 0 disables the gate for benchmarking the declined regime; an unparseable value falls back to the default. See sparse-paged-decode.md.
MLXCEL_SDPA_VECTOR_LARGE_D	0/false/off/no disable; any other value or unset enables	on	CUDA only. Gates whether the CUDA supports_sdpa_vector check accepts head_dim 256/288 (gemma family, qwen3.5/3.6, baichuan-m1, paligemma2), routing their decode to the fused sdpa_vector kernels instead of the materializing SDPA fallback (issue #675). Disabling restores the prior fallback with no rebuild; used for the A/B in benchmarks/cuda_gb10_sdpav_675_2026-07-06.csv.
MLXCEL_PIPELINE_GRANULARITY	off, layer, block:N	off	Inserts layer-boundary async-eval hints for pipeline experiments.
MLXCEL_FUSED_MOE	0/false/off/no disable; any other value or unset enables	on	Fused single-token decode-MoE kernel (#268), on by default since #282 (Metal) and #319 (CUDA, via mx.fast.cuda_kernel); validated on M1 Ultra, M5, and GB10. Set to 0 to force the proven gather_qmm/SwitchGLU path. Active for qwen3_moe, qwen3_next, dots.llm1, gemma4, qwen2_moe, mixtral, phimoe, lfm2, qwen3_vl_moe, and olmoe decode. Byte-identical greedy output is checkpoint-dependent and was never a general property (#1045): it holds on qwen3-30b-a3b and not on Klear. This is not a defect, since the kernel measures roughly 6x closer to an all-f32 ground truth than gather_qmm on both, but gather_qmm is what mlx-lm mirrors, so set this to 0 when reference-diffing a new MoE port.
MLXCEL_FUSED_MOE_SGY	1-32	8	Simdgroups (Metal) / warps-per-block (CUDA) per threadgroup for the fused decode-MoE kernel; tune per hardware.
MLXCEL_FUSED_MOE_MAX_DFF	positive int	4096 (Metal) / 8192 (CUDA)	Expert-intermediate (Dff) upper bound for the fused path; above it the caller falls back to gather_qmm. The fused path wins only while gather_qmm underutilizes the GPU (small experts), so the break-even is backend-dependent and the default is chosen from the live backend: 4096 on Metal (M1 Ultra tuning) and 8192 on CUDA (GB10 re-measured under MLX pin e9463bb, #626; fused wins through Dff 6400 and is break-even at 8192). An explicit value overrides the default on both backends: lower it to force gather_qmm sooner, raise it (e.g. 20000) to force the fused kernel on larger experts such as mixtral (Dff 14336, where it is a slight net loss).
MLXCEL_FUSED_MOE_RELU2	presence enables	off	Enables the squared-ReLU fused routed-expert path for nemotron-class experts. It is performance-neutral on nemotron-h even though MoE is its largest block because it replaces only the already-efficient routed fc1/fc2 GEMVs; the router, always-on shared expert, and combine remain outside this kernel. Kept for a future squared-ReLU model whose routed-expert slice dominates.
MLXCEL_FUSED_MOE_PARITY_CHECK	unset or 0 disables; nonzero integer enables	off	In-situ per-call parity/determinism probe for the fused decode-MoE kernel, added for the #886 corruption triage. Re-runs the fused kernel pair on identical inputs (bitwise-determinism check), computes the gather_qmm production-fallback reference, and computes an all-f32 dequantize-and-matmul ground truth, then logs to stderr any call that is non-deterministic or whose fused-vs-reference, fused-vs-ground-truth, or reference-vs-ground-truth deviation exceeds MLXCEL_FUSED_MOE_PARITY_THRESHOLD (normalized RMS). Heavy: adds several eager evals per MoE call. Diagnostic only, kept for future corruption triage; never enable in production.
MLXCEL_FUSED_MOE_PARITY_THRESHOLD	float	0.05	Normalized-RMS deviation threshold for the MLXCEL_FUSED_MOE_PARITY_CHECK probe.
MLXCEL_GATHER_QMM_GROUPED	0 disables; any other value or unset enables	on	CUDA only. Gates the sorted MoE prefill fast path in GatherQMM::eval_gpu (issue #629). When the sorted-indices M == 1 prefill contract holds (right-sorted, transpose, one activation row pre-gathered per (token, expert) pair, non-nvfp4, float activation dtype, E <= 1024) and the batch is large enough to amortize (see MLXCEL_GATHER_QMM_GROUPED_MIN_ROWS), the expert weight stack is dequantized once and routed through cutlass_grouped_gemm_unaligned instead of one 1-row qmm_sm80 GEMM per (token, expert) pair. Fixes a 5-10x CUDA MoE prefill collapse relative to Metal M1 Ultra; measured 3.6-40x prefill speedup across mixtral-8x7b, phi-3.5-moe, llama-4-scout, minimax-m2, gpt-oss-20b, and solar-open-100b on GB10, with no decode-path regression. Set to 0 to force the legacy per-row dispatch (A/B, rollback). See docs/benchmark_results/moe-prefill-grouped-gemm-gb10-2026-07-10.md.
MLXCEL_GATHER_QMM_GROUPED_MIN_ROWS	positive integer	8	CUDA only. Amortization threshold for MLXCEL_GATHER_QMM_GROUPED: the fast path activates only once the sorted batch size B reaches min_rows * E (E = expert count), the point past which the one-time expert dequant traffic is cheaper than the legacy per-row re-reads. Lower it to trigger the fast path on smaller sorted batches (e.g. high-concurrency batched decode crossing the same sort gate); raise it to keep more traffic on the legacy path for tuning or rollback.
MLXCEL_AUTOTUNE	1/true/on/yes enable tuning; cache/read/readonly/read-only consume cached tactics only; 0/false/off/no or unset disable	off	Shape-bucketed kernel autotuner (issue #906). Off by default and fully inert when off: no tactic-cache reads, no profiling, no filesystem access on the decode path, so cold-start latency and steady-state behavior are exactly what they were before #906. Set to cache to consume tactics an earlier mlxcel tune wrote without ever profiling at runtime, or to 1 to additionally profile the first use of a shape bucket that has no cached entry. Tuned tactics are keyed by (op, kernel/runner identity, power-of-two shape bucket, dtype) and stored one JSON file per key at ${MLXCEL_CACHE_DIR:-$HOME/.cache/mlxcel}/autotune/<key-hash>.json, carrying the mlxcel version and the pinned MLX commit they were measured under; a mismatch in either discards the entry and re-tunes rather than applying a stale config, and a corrupt or unreadable file is warned about and ignored, never fatal. A shape outside the tuned matrix warns once and uses the op's default. Selection is min-latency over a median-of-N timing per candidate, where N scales with the measured per-launch cost (at least 5, and many more for a cheap launch, which is where run-to-run noise concentrates); candidates are sampled round-robin so drift over the sweep moves all of them together rather than favoring whichever was measured first. A candidate must beat the op's own default by more than 2% and by more than the two measurements' combined relative spread (scaled median absolute deviation over the repetitions), and two candidates that cannot be told apart collapse to whichever sits nearer the default. A win inside the host's noise floor is therefore not a win, so a noisy host converges back to the default instead of flapping; the spread and the threshold that was applied are recorded in the cache entry and printed by mlxcel tune. An explicitly-set variable for a tuned knob (MLXCEL_PAGED_DECODE_SPLITS, MLXCEL_QMM_TILE_M, MLXCEL_QMV_MULTIROW, MLXCEL_QMV_MULTIROW_MAX_ROWS) always wins over a tuned value and suppresses tuning for that knob. Only the v1 paged-decode NumSplits op is tuned lazily at runtime; the process-wide CUDA kernel knobs are tuned only by the offline mlxcel tune subcommand and applied at process start.
MLXCEL_PAGED_DECODE_SPLITS	positive integer	unset (memory-budget ceiling)	Pins the NumSplits token-split launch shape of the v1 fused paged-attention decode kernel (issue #906), overriding both the built-in ceiling and any autotuned value. NumSplits is the number of SIMD groups per threadgroup that sweep strided token stripes; before #906 it was always the largest value the tg_acc[NumSplits * Dim] threadgroup-memory budget and the 1024-thread cap allow, which is a feasibility bound rather than a performance optimum (it ignores context length and batch). Values are clamped into [1, ceiling] by the C++ launcher, so an over-large value is safe; a non-positive or unparseable value is warned about and ignored. Each value is a distinct JIT specialization, so switching costs a one-time kernel compilation. Use for A/B against a tuned choice or to reproduce a recorded benchmark configuration.
MLXCEL_QMV_MULTIROW_MAX_ROWS	1-8	8	CUDA only. Row-window ceiling for the multirow qmv path (issue #906), narrowing the 2 <= M*B <= 8 dispatch window that MLXCEL_QMV_MULTIROW (issue #725) introduced. The crossover past which the small-M qmm shape takes over is a per-hardware property, not a constant (docs/CONTINUOUS_BATCHING.md documents a regression past 7 rows on GB10), so the autotuner tunes it and mlxcel tune --op qmv-multirow publishes the winner here at process start. The window can only be narrowed, never widened: 8 is the widest accumulator array the multirow kernel instantiates, so values above 8 (and unparseable values) fall back to 8. Narrowing it no longer costs registers that the launch does not use: since the row-blocking change the accumulator width is dispatched on the actual row count (2, 4 or 8), so a 4-row launch takes 68 to 71 registers rather than the 95 to 109 an 8-row instantiation needs on sm_70. A ceiling of 1 disables the multirow path entirely, which is equivalent to MLXCEL_QMV_MULTIROW=0. An explicitly-set value always wins over a tuned one. Unvalidated: no CUDA host was available when this landed, so tune and compare against docs/benchmark_results/qmv-multirow-gb10-2026-07-11.md before trusting a tuned value.
MLXCEL_BENCH_LLC_BYTES	positive integer (bytes)	detected estimate	Overrides the last-level-cache size the cold-L2 benchmark mode uses to size its input rotation (issue #906). The default is an Apple SLC estimate by device family (macOS exposes no SLC size through sysctl), deliberately biased high because under-estimating silently reintroduces the warm-cache bias the mode exists to remove; on CUDA hosts the estimate falls back to an 8 MiB floor because reading cudaDeviceProp::l2CacheSize needs an FFI helper that does not exist yet, so set this to the device's real L2 size there. Only read by benchmark harnesses (examples/page_gather_microbench.rs --cold-l2); it has no effect on inference. See docs/benchmarks.md.
MLXCEL_QMV_MULTIROW	0 disables; any other value or unset enables	on	CUDA only. Gates the weight-amortizing multirow qmv path (issue #725). With broadcast weights and 2 <= M*B <= 8 input rows (batched decode with B in [2,8), speculative-verify [1, K] forwards), one warp applies each dequantized weight tile to every input row instead of launching one weight-rereading block column per row, so weight DRAM traffic is O(1) in the row count. Per-row outputs are bit-identical to the stock per-row launches (pinned by qmv_multirow_matches_per_row_qmv_bitwise); classic M*B == 1 decode is untouched. On GB10 this flips B=1 MTP speculative decoding from 0.52-0.77x to 1.31-1.46x and lifts 4-client aggregate serving decode from ~50 to ~74 tok/s on llama-3.1-8b-4bit. Set to 0 to force the stock per-row dispatch (A/B, rollback). See docs/benchmark_results/qmv-multirow-gb10-2026-07-11.md.
MLXCEL_QMV_ROWS_PER_WARP	1, 2 or 4; any other value (or unset) uses the per-architecture default	2 on sm_70, 1 elsewhere	CUDA only. Output rows each warp covers in qmv, for classic M*B == 1 decode. The stock kernel gives one warp one output row and has that warp read the whole activation vector, so a launch moves n * k * 2 bytes of activations against n * k / 2 bytes of 4-bit weights. Measured on a V100 over a controlled 4-bit against 8-bit pair of gemma-4-12B-it with identical launch counts, that activation term is 84% of a 4-bit launch, which is why decode below Ampere is not weight-bandwidth bound. R rows per warp load the activation tile once and apply it to all R weight rows, dividing the activation traffic by R while leaving weight traffic alone; the per-row arithmetic is unchanged, so output is byte-identical to R = 1 at temperature 0. The cost is that the same output rows are covered by R times fewer warps and the accumulator array grows to R * elems_per_thread. On a V100 at 4 bits, R = 2 moves decode from 44.7 to 60.1 tok/s on gemma-4-12B-it-4bit under the MLXCEL_QMV_MIN_BLOCKS default of 3, and from 20.6 to 24.9 on qwen3.8-27B-4bit; R = 4 gives it all back on the 12B (96 to 142 registers, 1 to 2 resident blocks per SM against the stock kernel's 4 to 5) while costing nothing on the wider 27B. Turing defaults to 1 because no sm_75 device was available to measure and the R = 4 reversal shows the tradeoff does not extrapolate. 1 routes to the stock kernel with the stock grid, so it is the rollback. Read once per process.
MLXCEL_QMV_MIN_BLOCKS	1, 3 or 4; any other value (or unset) uses the per-architecture default	3 on sm_70, 1 elsewhere	CUDA only. Resident-block floor handed to __launch_bounds__ on the row-blocked qmv kernel, which is how ptxas is told to cap the register count: for a 256-thread block, 3 caps at 85 registers per thread and 4 caps at 64. This is the knob, not the row count. Measured on a V100 at 4 bits on gemma-4-12B-it, decode as the slope between -n 60 and -n 200: R = 2 unconstrained sits at 92 to 95 registers and 2 blocks per SM for 56.0 tok/s, the same R = 2 under a floor of 3 sits at 80 registers and 3 blocks for 60.1, and under a floor of 4 it sits at 64 registers and falls back to 54.9, so the budget can be set too tight as well as too loose. Nothing spills at any floor (LOCAL is 0 in every instantiation). R = 3 under the same floor of 3 has identical registers and residency and moves 17% fewer memory transactions per output row, and it still lands at 58.8, which is why the floor rather than the row count carries the default. 1 is the unconstrained case and is the rollback. Read once per process.
MLXCEL_SSM_CUDA_KERNEL	0 disables; any other value or unset enables	on	CUDA only. Gates the fused single-token SSM decode kernel port (issue #631). The Metal-only ssm_update_kernel (one launch replacing the ~55-op SSD scan graph per SSM layer) now has a mx.fast.cuda_kernel port, selected automatically on the CUDA backend for the hybrid SSM models (granite-4.0-h, falcon-h1, plamo-2, nemotron-h). Fixes the 0.29-0.36x hybrid-SSM decode ratio vs Metal M1 Ultra; measured 2.6-4.5x decode speedup on the granite/falcon family on GB10, greedy parity byte-identical, pure-mamba2 control unchanged. Set to 0 to force the graph path (A/B, rollback). See docs/benchmark_results/hybrid-ssm-decode-cuda-kernel-gb10-2026-07-10.md.
MLXCEL_FUSED_QK_NORM	1/true/on/yes enable; any other value or unset disables	off (opt-in)	Fused single-token QKV projection + Q/K RMSNorm + RoPE kernel (#326) for Qwen3 and Qwen3-MoE decode. Opt-in: set to 1 to enable. Matches the graph path within RMS < 5e-3 (the reduction is over the transpose-invariant head_dim axis), but greedy temp-0 is not byte-identical over long generation; on CUDA the graph path is itself non-deterministic run-to-run from GPU FP-reduction order, while the fused path is deterministic, so its output stays inside the graph baseline's own envelope. The kernel cuts Rust/C++ FFI crossings rather than MLX op count, so it does not speed up the GPU/bandwidth-bound decode loop: on M1 Ultra it measured 1 to 3.4% slower (qwen3-0.6b 275 vs 284, qwen3-8b 82.3 vs 83.2 tok/s); on GB10/CUDA (SM 12.1) it is also slower (qwen3-0.6b 0.96x, qwen3-8b ~1.0x, qwen3-30b-a3b 0.92x fused/graph; see docs/benchmark_results/fused-qk-norm-decode-gb10.md), so there is no per-backend win. Ships as a reusable shared primitive for the deferred QK-norm families and stays opt-in (default off) on every measured backend (M1 Ultra, M5 Max, GB10/CUDA), mirroring the opt-in MLXCEL_FUSED_MOE_RELU2. Active only when l == 1 (decode) and weights are quantized.
MLXCEL_FUSED_ADD_RMSNORM	0/false/off/no disable; 1/true/on/yes enable; unset or unrecognised keeps the default	off (opt-in)	Fused residual-add + RMSNorm decode kernel (#905). Computes new_residual = x + residual and normed = rms_norm(new_residual) * (weight_bias + weight) in one dispatch instead of an elementwise Add followed by fast::rms_norm, removing one kernel launch and one full-width intermediate at every pre-norm residual join. Adopted by the Llama3-family block (which also serves Qwen2/Qwen2.5 and every VLM whose text backbone is Llama3Model/Qwen2Model) and by Gemma, where the (1 + w) convention rides in as a scalar weight_bias = 1.0 folded in the weight's own dtype rather than as a second weight tensor. Numerics are pinned to MLX's own rms_norm kernel (fp32 accumulation over the dtype-rounded sum, precise::rsqrt, the x * inv_mean product rounded to the activation dtype before the weight multiply), so the fused and unfused paths agree to within a small multiple of one ulp; greedy temp-0 is not guaranteed byte-identical over long generation, because a near-tie argmax can land on either side of that difference. Set to 0 to force the add + fast_rms_norm pair (A/B, rollback). The compiled-in default lives in FUSED_ADD_RMSNORM_DEFAULT in src/lib/mlxcel-core/src/layers.rs. Read once at first use and cached for the process lifetime.
MLXCEL_FUSED_ROPE_APPEND	0/false/off/no disable; 1/true/on/yes enable; unset or unrecognised keeps the default	off (opt-in)	Fused q/k RoPE + KV-append-layout decode kernel (#905). Reads the row-contiguous fused-QKV projection output whole, applies rotary embedding to the q and k blocks, and emits q, k and v already in their consumers' layouts, replacing three trailing-axis slices, three reshape/transpose pairs and two fast::rope calls with one dispatch. The rotation is transcribed from MLX's own rope.metal, so it agrees with fast::rope including the traditional convention and partial rope_dims. The trig call is the one thing chosen per backend rather than transcribed: each body calls whatever MLX's own RoPE kernel for that backend calls, so the Metal body uses metal::fast::cos/sin after rope.metal and the CUDA body uses libdevice cosf/sinf after rope.cu. Using the CUDA fast intrinsics __cosf/__sinf there instead cost about 1.2e-2 rad of applied phase at absolute position 131071 and failed the parity suite at the two large-offset cases (issue #1049). Adopted by the Llama3-family dense decode path, except for a checkpoint whose rope_scaling block selects a frequency table: the kernel derives its frequencies from rope_base and has no parameter for a table, so such a model takes the graph path (#1355), and a one-time notice on stderr says so when this variable asked for the kernel. traditional and the position scale are real parameters and are honored, so neither triggers that notice for this kernel. The paged block-pool destination layout the kernel also supports is implemented and tested but not wired, because the batched paged decode path belongs to issue #899. The kernel emits the append payload rather than writing the cache slab, because MLX custom-kernel outputs are always freshly allocated while the existing slice_update append donates the slab buffer, so an in-kernel write would trade an O(1) donated update for an O(capacity) copy. Set to 0 to force the reshape/transpose/fast_rope graph (A/B, rollback). The compiled-in default lives in FUSED_ROPE_APPEND_DEFAULT in src/lib/mlxcel-core/src/layers.rs. Read once at first use and cached for the process lifetime.
MLXCEL_COMPILED_QGELU_MLP	0/false/off/no disable; any other value or unset enables	on	Compiles the affine-quantized GeGLU MLP (gate/up/gelu/down) into one mx::compile graph so the tanh-approx GELU's ~14 element-wise ops per layer collapse into a single fused Compiled primitive. Covers the Gemma family (gemma/gemma2/gemma3/gemma4). The group_size=64/bits=4 case was already compiled on every shape and is unchanged; other affine quantizations (notably the group_size=64/bits=8 MLP weights in Gemma 4 mixed-precision checkpoints, issue #680) are compiled only on the single-token decode call, because compiling the 8-bit prefill GEMM measured 8-9% slower and +0.2-0.7 GB peak on GB10 (the shapeless fused graph forces a decode-oriented qmm kernel onto the large prefill matmul). GB10 gemma-4-12b decode is weight-bandwidth-bound (~94% GPU-busy), so this primitive-count cut is throughput-neutral there (measured <1% from a ~530-primitive/step reduction); it still removes real CPU dispatch and helps op-count-bound backends. Set to 0 to force the op-at-a-time fallback (A/B + rollback). It also gates the NVFP4 scaled fused MLP variant (MLXCEL_DISABLE_FUSED_GLOBAL_SCALE), forcing its eager fold when set to 0.
MLXCEL_DISABLE_FUSED_GLOBAL_SCALE	1/true/on/yes (case-insensitive) disable; unset or any other value keeps the fold	off (fold on)	Rolls back the NVFP4 global-scale fold for Gemma 4 (issues #698/#705). By default the fused MLP and per-layer-input-gate C++ paths fold each per-projection weight_scale_2 sidecar (from the direct ModelOpt transcode, issue #693/#697) into the fused kernel at the mathematically correct points: the gate scale before the GeGLU activation, the up scale on the up product, and the down scale on the fused output, each reproducing apply_global_scale byte-for-byte. Native NVFP4 prefill now uses a shape-specific scaled MLP graph, and standalone UnifiedLinear sidecar projections can apply qmm + global scale + dense bias through one C++ helper. When set, sidecar-carrying paths fall back to the op-at-a-time UnifiedLinear::forward scalar application (the pre-#698 bypass). Greedy temp-0 decode is token-identical across the two paths on gemma-4-31b-it-nvfp4; the fold removes element-wise dispatches on the gemma-4 path where CUDA graphs are disabled (#688). See docs/benchmark_results/nvfp4-direct-transcode-gb10-2026-07-08.md and docs/benchmark_results/nvfp4-native-prefill-m1ultra-2026-07-09.md.
MLXCEL_FUSED_XIELU	0/false/off/no disable; any other value or unset enables	on	Fused single-launch Metal xIELU kernel for the Apertus MLP activation (#409), on by default since the M5 Max validation. MLP::forward routes through one Metal dispatch covering the ~11 elementwise ops in apertus_xielu (square, minimum, expm1, where, and neighbors) instead of the per-op graph. Greedy temp-0 decode is byte-identical to the elementwise path on Apple Silicon: every intermediate stays in the input dtype (bf16) and the kernel reproduces MLX's expm1f exactly. Measured decode speedup on M1 Ultra (+2.7%, Apertus-8B 83.4 to 85.7 tok/s) and M5 Max (+1.9%, 112.0 to 114.2 tok/s), with no regression. Set to 0 to force the elementwise path. On non-Metal back-ends the FFI falls back to an equivalent elementwise graph, so the flag is safe to set everywhere. Apertus only; no other model family is affected.
MLXCEL_SAMPLING_GUMBEL	0/false/off/no disable (case-insensitive); any other value or unset enables	on	Softmax-free Gumbel-max categorical sampling kernel (issue #900), on by default on Metal and CUDA. On the no-filter stochastic path (temperature > 0 with no top-k, top-p, or min-p) the fused sampler replaces mlx::core::random::categorical with one index-carrying max reduction over the vocabulary: adding i.i.d. Gumbel(0,1) noise to logits / temperature and taking the argmax draws exactly from softmax(logits / temperature), so the normalization pass over 32K-152K entries disappears. Noise comes from a stateless in-kernel Philox-4x32-10 counter keyed on (row, element), so the sampled id is a pure function of the launch key and the logits and does not move when the launcher splits a row across threadgroups to keep a small batch wide. One launch covers the whole [B, vocab] batch. This changes the RNG stream: at an equal seed the sampled tokens differ from the pre-#900 categorical streams, because the two consume the shared MLX key sequence differently. The distribution is identical (chi-square goodness of fit against exact softmax is a committed test), and greedy (temperature == 0) output is byte-identical either way. Set to a falsy value to restore categorical exactly, which is the way to reproduce a token stream recorded before this landed. Filtered configurations (top-k / top-p / min-p) never reach the kernel and are bit-identical with the switch either way; token bias is an upstream -inf logit mask that the kernel reproduces exactly, and XTC-active requests never reach this kernel (XTC lives inside the stock filter chain). Read once per process, so set it before starting mlxcel or mlxcel-server. Inert on CPU-only builds and under MLXCEL_DEVICE=cpu. A/B with examples/gumbel_sampling_microbench.rs.
MLXCEL_SAMPLING_REJECTION	0/false/off/no disable (case-insensitive); any other value or unset enables	on	Sorting-free top-k / top-p / min-p sampling by dual-pivot rejection (issue #901), on by default on Metal and CUDA. On the filtered stochastic path the fused sampler replaces the argpartition top-k, the argsort + cumsum top-p and the min-p mask, plus the trailing random::categorical, with the softmax passes and one custom kernel; since the untempered-chain fix the kernel reads two probability rows (the untempered row resolves every filter, the tempered row weights the draw), which coincide at temperature 1. Routing is restricted to the configurations measured faster. The kernel replaces a sort, so it is routed only where the stock chain sorts, which is when top-p is active; and when top-k is active as well, only at vocabularies up to 32768. Measured on M1 Ultra over three repetitions of vocab {32K, 64K, 152K} x batch {1, 4, 8}: top-p alone 1.28x to 2.35x at every cell, top-k alone 0.31x to 0.97x at every cell, min-p alone 0.47x to 0.88x at 152K, top-k with top-p 1.27x to 1.64x at 32K but 0.71x to 0.83x at 152K. The mechanism is the kernel's round count, which the microbenchmark reports: top-p accepts in one or two rounds because a probability-weighted draw lands in the high-mass head, while top-k needs the draw to land in the top forty of the whole vocabulary and takes two to seven rounds, growing with the vocabulary. Every round is another full-row sweep, and the loop runs on one threadgroup per row because the shrinking interval is carried across sweeps and MLX custom kernels have no grid-wide barrier, so the kernel pays single-core bandwidth per round while argpartition and the stock min-p mask run across the whole GPU. top-k alone, min-p alone, top-k with min-p, and top-k with top-p above vocab 32768 therefore stay on the stock chain and are bit-identical to before; the decline is announced once at INFO with the numbers behind it. Vocab 65536 has not been measured for the top-k plus top-p combination and is excluded rather than interpolated. All three filters are threshold filters on the probability value, so the filtered support is {p > low} for a single scalar per row; the kernel finds that scalar without sorting, by rejection sampling on a shrinking interval. Each round draws a candidate from the current proposal with a fixed-order block scan, takes the candidate's probability as one pivot and the midpoint of the bracket as the other, reduces (count, mass) above both in one vocabulary sweep, and accepts the candidate as soon as it passes every active filter test. The bisection pivot is the midpoint of the IEEE-754 bit patterns, not the arithmetic mean: pure integer arithmetic that a fast-math backend with flush-to-zero cannot corrupt for subnormal probabilities, and it isolates a single float in at most 32 rounds at any magnitude. min-p is folded into the initial interval and costs no round at all. One launch covers the whole [B, vocab] batch, and the kernel reads {top_k, top_p, min_p} per row so rows with different values need no second launch. This changes the RNG stream: at an equal seed the sampled tokens differ from the pre-#901 streams, because the kernel consumes the shared MLX key sequence differently from random::categorical. One semantic difference is intentional. When top-k AND top-p are both active, the stock chain masks to the top-k set and renormalises before applying top-p, so its mass target is top_p * Z_k with Z_k the top-k mass; the kernel applies both tests to the untruncated distribution, so its target is top_p * total and its support is a superset. The two agree unless some token's exclusive cumulative mass falls in (top_p * Z_k, top_p * total]. Every other configuration (top-k alone, top-p alone, min-p alone, top-k+min-p, top-p+min-p) is exact, and support equality against the argpartition mask including ties, plus chi-square goodness of fit against the renormalised truncated distribution, are committed tests. Greedy (temperature == 0 or top_k == 1) is byte-identical either way, and the no-filter path is unaffected (it belongs to MLXCEL_SAMPLING_GUMBEL). The routed path never synchronizes. Both decode drivers are software pipelines that build step n+1 and async_eval it before reading step n, so a sampler that evaluates anything internally drains the queue inside the caller's build phase and collapses the pipeline. The first cut of this kernel read the per-row converged flags back to host inside fused_sample and measured 1.7x slower end to end on Qwen3-0.6B, against an op-level matrix that scored the same configuration 1.14x to 1.17x faster. The production path therefore evaluates nothing: the round cap is unreachable by construction (the bit-space bracket bounds the loop at 31 rounds), and the converged flags are stashed and inspected on a later call once array::is_available() reports the launch has landed, which in a decode loop is by the next token. A row that somehow exhausted the cap returns its own argmax, which is always inside the filtered support, so the degradation would be one greedy draw rather than an invalid token; the event is counted in mlxcel_core::rejection_cap_overflow_rows() and announced once at INFO. Which path a run took is announced at INFO, once per distinct outcome kind, so a benchmark arm can be proven from a log. Set to a falsy value to restore the argpartition chain exactly, which is the way to reproduce a token stream recorded before this landed. Read once per process, so set it before starting mlxcel or mlxcel-server. Inert on CPU-only builds and under MLXCEL_DEVICE=cpu. A/B with examples/rejection_sampling_microbench.rs.
MLXCEL_CUDA_F16_NORMALIZE	1/true/on/yes enable; 0/false/off/no disable	on below sm_80, off (opt-in) at sm_80 and later	CUDA only. Load-time bf16 -> f16 normalization of weights for the single-dtype decode graph (issue #636). The default is split by compute capability because the hardware is. Ampere and later (sm_80+): off, unchanged from #636. bf16 has native ALUs there, the merged patches-cuda/dtype.cpp promotion patch already yields a 0-AsType single-dtype bf16 decode graph, and f16 offers no measured throughput gain (qwen2.5-0.5b-bf16 decode 208 vs 207 tok/s on GB10) while narrowing dynamic range. Volta and Turing (sm_70, sm_75): on, and quantized checkpoints are included. Those parts have no bf16 ALU, no bf16 tensor-core MMA atom, and no cuBLAS bf16 GEMM, so every bf16 operand is converted before it can execute. Measured on a Tesla V100 with qwen3.8-27B-4bit: prefill 85.3 s to 4.8 s and decode 118.1 to 49.2 ms per token, with qmm_naive spending 199.1 s against 11.4 s over an identical 994 launches. On a quantized checkpoint the packed planes stay u32 and only the bf16 side-data moves, which is why quantized models are in scope here and out of scope above. Set to 0 to opt out on pre-Ampere and keep bf16. The f16-fragile exception differs by architecture. At sm_80 and later it is the full conservative list from #732 (gemma, cohere/command-r, apertus, gpt-oss, and any config with a nonzero softcap or logit_scale). Below sm_80 every entry on that list was measured on a V100 and only apertus stays excluded: its xIELU squares its input, and f16 produces NaN at all 508 scored positions from the first token where bf16 scores 24.27. Gemma (+0.09% perplexity over 10208 tokens, decode 6.12 to 9.79 tok/s), gpt-oss (token-weighted -0.02%, with the two window sizes disagreeing in sign) and Cohere (which ships F16 weights already and so runs in f16 regardless, scoring 20.25 and 13.15) all measured clean, as did the generic softcap and logit_scale triggers, which is what a cap bounding a value rather than growing it predicts. Metal / Apple Silicon is governed by the separate always-on policy and is untouched.
MLXCEL_CUDA_F16_FRAGILE	1/true/on/yes enable; unset keeps the exclusion	off (opt-in)	CUDA below sm_80 only. Converts a family on the f16-fragile exception list anyway. That list is inherited from #732, which adopted it for Ampere and later where f16 gave, in that commit's words, "no AsType reduction and no throughput gain": with no upside, excluding a family on a suspicion costs nothing, and that PR's test plan ran no fragile family in f16. Below sm_80 the list is down to apertus, which is the one entry measurement supported, so this flag now means converting a family that produced NaN at every scored position on this hardware. It remains for a checkpoint whose activations you have reason to believe differ, and for taking the measurement on a family this list does not yet know.
MLXCEL_KEEP_BF16	set to any value to enable; unset disables	off	Measurement instrument, not a supported configuration. Skips load-time bf16 -> f16 conversion for every checkpoint, on every backend, so a bf16 model loads uniformly bf16. It exists to A/B a dtype policy against the mixture it replaces, and it is what established the Gemma3n result below. Gemma3n used to keep only its language MLP bf16 and convert the rest to f16, leaving one checkpoint holding both dtypes; MLX promotes f16 with bf16 to f32, so every boundary between the kept MLP and its f16 neighbours promoted and the decode path read promoted weights the whole way down. Loading uniformly bf16 measured, on M5 Max with gemma3n-e4b-bf16 at a 273-token prompt with an image and 128 generated, decode 38.02 to 47.48 tok/s and prefill 2143 to 2371 tok/s, against mlx-vlm 0.6.17 at 48.72 and 1207: decode moves from 78% of the reference to 97% while prefill stays about 1.95x ahead. Gemma3n now loads uniformly bf16 by default (see gemma3n_bf16_key), so this flag no longer changes that family. It is not evidence that bf16 beats f16. The control run is the point: on families with no forced-bf16 subset the two dtypes land within noise of each other on the same hardware, qwen2.5-0.5b-bf16 at 401.8 against 397.3 tok/s and llama-3.1-8b-bf16 at 33.07 against 33.02. What costs is the mixture, not the dtype, so every other family keeps converting to f16 and setting this flag for them buys nothing while narrowing nothing. Read per load.

Block-diffusion diagnostic variables

Variable	Values	Default	Purpose
MLXCEL_DIFFUSION_DEBUG_CANVAS=1	1 enables	off	Diagnostic. Replaces all DiffusionGemma canvas random-noise initialization with a fixed deterministic pattern ((i+1)*7919 + k*104729) % vocab_size) and prints DIFFUSION_COMMIT block=<n> ids=... per committed block. Intended for cross-implementation parity testing against the mlx-vlm Python reference at temperature 0. Output is not suitable for normal generation.

Logging, profiling, and capture variables

Most of these switches force synchronization or extra graph work and will change throughput measurements. Use them for diagnosis, not capacity planning.

Variable	Values	Default	Purpose
MLXCEL_TRACE_ARCH	presence enables	off	Prints the CUDA architecture picture once per process, to stderr, at device init (issue #1537): the running GPU's compute capability, the MLX_CUDA_ARCHITECTURES list this binary was compiled for, and whether the device is served by a precompiled cubin or by JIT-compiled PTX. On a CUDA build it also prints the quantized-matmul path the dispatcher picks on its first call (qmm_sm90, qmm_sm80, qmm_naive, qmv, or fp_qmv), which is the most architecture-dependent choice in decode and which MLX exposes no other way. Diagnostic only, no synchronization and no extra graph work; on a Metal or CPU-only build it reports that no compute capability is available. Note that an architecture mismatch is reported without this variable: a binary whose compiled architectures do not cover the host GPU refuses to start with a named error naming both, which MLXCEL_DEVICE=cpu bypasses.
MLXCEL_TRACE_DTYPE	presence enables	off	Prints selected tensor dtypes/shapes during generation.
MLXCEL_TRACE_ASTYPE	presence enables; 2/break adds breakdown	off	Prints the AsType (dtype-conversion) node count in the first decode step's graph (the single-dtype decode-graph metric, issue #636). Set to 2 or a value containing break to also dump the per src->dst dtype breakdown. Graph traversal only, no extra eval; zero cost when unset. Do not combine with MLXCEL_TRACE_DTYPE, which pre-evaluates the logits and collapses the graph before the count.
MLXCEL_DISABLE_CACHE_WARMUP	1 disables (exact match only; any other value, including true/on, or unset leaves the feature ON)	off, i.e. the warm-up is enabled	Kill switch for the background prompt-cache warm-up (issue #1144), which extends a conversation's history-boundary snapshot to cover the previous assistant reply so the next turn prefills only the new user message. Deliberately separate from MLXCEL_DISABLE_BOUNDARY_SNAPSHOT: the boundary snapshot is foreground work that buys the next turn a hit, while the warm-up is speculative background work that makes that hit bigger, and an operator may reasonably want the first without the second. Warm-ups run only when the scheduler is fully idle (no decode batch, no queued or parked prefill), so they never delay a foreground request. Read once through a OnceLock, so setting it after the server has started has no effect.
MLXCEL_DISABLE_BOUNDARY_SNAPSHOT	1 disables (exact match only; any other value, including true/on, or unset leaves the feature ON)	off, i.e. the boundary snapshot is enabled	Kill switch for the history-boundary prompt-cache snapshot (issue #1143). That feature is what makes multi-turn prompt caching hit at all for snapshot-only families (Gemma 4, Qwen 3.5, and the SSM hybrids), so it is on by default. It costs a second chat-template render and a second tokenization on the request-dispatch thread plus an extra graph launch and a model-state copy on the foreground prefill, which a deployment serving only single-turn traffic pays for reuse it never claims; set this to 1 to restore the previous request path end to end. Like MLXCEL_APC_TRACE the value is a literal string match read once through a OnceLock, so setting or changing it after the server has started has no effect. Note that turning this off also makes the prefill shape identical to the pre-#1143 one, which can shift a greedy near-tie token relative to a run with it on.
MLXCEL_APC_TRACE	1 enables (exact match only; any other value, including true/on, or unset leaves it off)	off	Opt-in prompt-cache trace logging (issue #774). Unlike the presence-based switches on this page, the check is a literal string match, not a truthy/falsy parse. The value is read once, effectively at first use, via a OnceLock, so setting or changing it after the server process has started has no effect. When enabled, the scheduler emits one tracing::info! line per prompt-cache store, adopt, and reject event, each carrying the sequence id and the relevant token/prefix length (matched length for store and adopt, context length for reject), with reject lines also carrying the specific PromptCacheRejectReason.
MLXCEL_FORCE_SYNC	presence enables	off	Forces synchronous decode evaluation. Also disables the server BatchScheduler's lookahead decode pipeline (issue #632), falling back to the pre-pipeline synchronous tick.
MLXCEL_PROFILE_PIPELINE	presence enables	off	Emits high-level generation pipeline timing.
MLXCEL_PROFILE_PIPELINE_DETAIL	presence enables	off	Emits one [PIPELINE_DETAIL] line per generation splitting the decode step into reshape, forward (host graph build), sample, async_eval (device work) and item_wait (host sync), in ms per token. Wired into both generate_streaming and generate_with_stats, so it fires under mlxcel-bench-decode as well as the server. This is the tool for a decode-throughput question: forward is roughly constant with model size while async_eval scales with it, so a small model whose ratio against mlx-lm is low is usually paying host build rather than device time. Costs a handful of Instant::now calls per token when enabled.
MLXCEL_PROFILE_TTFT	presence enables	off	Emits one [TTFT] line per --profile generation breaking the pre-first-token phase into setup, build (lazy graph construction, host only), sample, eval (the one blocking evaluation, which also materializes every weight, loads every JIT module and instantiates every CUDA graph) and post, plus the reported prefill and the unattributed residual. MLXCEL_PROFILE_PIPELINE covers only the decode loop, so this is the pre-first-token counterpart (issue #1545).
MLXCEL_PROFILE_BLOCKS	presence enables	off	Emits per-block/model-family timing where implemented.
MLXCEL_PROFILE_FORWARD	presence enables	off	Enables model-specific forward profiling where implemented.
MLXCEL_PROFILE_NEMOTRON_MOE	presence enables	off	Forces evaluation at the router, routed-expert, combine, and shared-expert boundaries inside fused_moe_forward, then emits one timing line per MoE call. Diagnostic only: the forced synchronization changes absolute throughput. On the experimental MLXCEL_FUSED_MOE_RELU2 path, routed_ms includes the score multiply folded into the down kernel and combine_ms is only the final top-k reduction.
MLXCEL_PROFILE_QWEN3_MOE_DETAIL	presence enables	off	Profiles Qwen3 MoE internals.
MLXCEL_PROFILE_MOE_INNER	presence enables	off	Profiles Gemma 4 MoE sub-operations.
MLXCEL_PROFILE_PER_LAYER	presence enables	off	Prints per-layer Gemma 4 timing.
MLXCEL_PROFILE_LAYER_BUILD	presence enables	off	Adds Gemma 4 layer-build timing.
MLXCEL_PROFILE_LAYER_SUBOPS	presence enables	off	Adds Gemma 4 per-suboperation timing.
MLXCEL_EXPORT_DECODE_DOT	file path	unset	Exports the first decode graph pair to DOT.
MLXCEL_CAPTURE_DECODE	path	unset	Captures one warmed decode token to a Metal GPU trace and exits; requires MTL_CAPTURE_ENABLED=1.
MLXCEL_METAL_CAPTURE_PATH	file path	unset	Starts a Metal capture around steady-state generation; requires MTL_CAPTURE_ENABLED=1.
MLXCEL_DEBUG_GEMMA4_LOAD	presence enables	off	Emits Gemma 4 safetensors loading diagnostics.
MLXCEL_NO_PRECISION_WARNING	presence suppresses	warning on	Suppresses the bf16-on-Apple-Silicon precision/performance note.

Test and CI variables

These are intended for the repository's own tests and automation rather than normal end-user operation.

Variable	Purpose
MLXCEL_CI_PP_MODEL	Model path used by the pipeline-parallel CI integration test.
MLXCEL_SKIP_HEAVY_TESTS	Skips selected heavy tests.
MLXCEL_BENCH_DATE	Metadata override for Turbo KV benchmark tests.
MLXCEL_BENCH_MACHINE	Metadata override for Turbo KV benchmark tests.
MLXCEL_ALLOW_PARALLEL_CUDA_TESTS	Downgrades the CUDA serialization guard to a warning. The guard fails when the mlxcel-core suite is about to run with more than one test thread under --features cuda, which aborts inside MLX partway through (#1048). Set it only to reproduce that abort, or when a narrow filter happens to match the guard's own name on a run that would have been safe.
MLXCEL_ALLOW_CONCURRENT_GPU_TESTS	Downgrades the GPU-exclusivity guard to a warning. The guard fails when a second mlxcel-core test binary is running, because two suites on one device corrupt each other and can report failures that do not exist (#1008).
MLXCEL_TEST_OP_PARITY_SEED	Operand seed for tests/metal_block_vs_chain_op_parity.rs. Defaults to 1165. The diagnostic compares one T = K call against K single-position calls byte for byte, and when a kernel pair differs by only a byte or two out of ten thousand, whether any byte differs is itself draw-dependent, so a single seed can report equal for a shape that is genuinely on a different kernel. Sweeping this separates that from a real equality. Test-only.
MLXCEL_MTP_DRAFT_PROFILE	1/true/yes/on to enable; unset or anything else to disable
MLXCEL_TEST_CHAIN_BLOCK	Verify-block width used by tests/qwen38_mtp_chain_parity.rs's block_verify_chain_matches_single_token_chain and adapter_chain_with_perfect_drafts_matches_classic_chain. Defaults to 3, the published drafter width; values below 2 are ignored. The byte-identity those tests pin is width-dependent, because MLX changes quantized-matmul kernel above get_qmv_batch_limit (12 for the 27B operand sizes on an M1 Ultra), so this is how that boundary gets measured on a given GPU without editing the source. Test-only.
MLXCEL_TEST_VIDEO	Turns the graceful skip in an ffmpeg-backed video test into a hard test failure. Only the exact string 1 enables it. The gated tests are the ffmpeg-backed ones in src/multimodal/video_tests.rs plus process_videos_pixel_values_match_input_color in src/vision/processors/gemma4_tests.rs, which decodes a synthetic clip through the same load_video entry point before running it through the Gemma 4 processor. All of them are #[ignore], so selecting them at all needs --include-ignored; this variable then asserts that the host can actually run them instead of skipping when ffmpeg/ffprobe is missing. make verify-test-video sets both, and nightly-verify.yml runs that target after installing ffmpeg. Test-only: no production code path reads it.
MLXCEL_REQUIRE_PINNED_CHECKPOINTS	Turns the graceful skip in a pinned-checkpoint contract test into a hard test failure. Only the exact string 1 enables it, which is narrower than the truthy parsing most variables on this page accept; unset leaves skips as skips. Currently gates pinned_post_tower_weight_roots_and_shapes_match_published_contract (src/vision/encoders/muse_glimmer_fusion_pinned_tests.rs) and pinned_weight_index_classifies_each_source_weight_once (src/loading/vlm_muse_glimmer_tests.rs), both of which validate the pinned Muse Glimmer checkpoint under models/mlx/muse-glimmer-30b against a published contract. Test-only: it has no effect outside cargo test and no production code path reads it.

This is a test-only gate: set it in a shell or CI job that runs cargo test, never in a runtime environment for mlxcel or mlxcel-server. It is meant to be enabled on a machine that owns the pinned checkpoint, because that is the only machine where the guarded tests do real work; everywhere else the checkpoint is absent by design, and skipping there is correct rather than a loss. On a checkpoint-owning machine, an unconditional skip silently disables the only coverage that checkpoint has, so leaving the gate off there hides exactly the failures it exists to catch.

To enable it for one run, set the variable ahead of the narrow test target, for example MLXCEL_REQUIRE_PINNED_CHECKPOINTS=1 cargo test --lib vision::encoders::muse_glimmer_fusion. On a machine that owns the pinned checkpoint permanently, export it in that machine's shell profile so every local test run enforces the gate without remembering to set it each time; nothing under .github/ or scripts/ references this variable today, because the checkpoint lives on a specific machine's local model store rather than in CI, and a repository change cannot export a variable into a maintainer's shell.

The gate has a second, already realized use beyond catching corruption: it turns an ambiguous pass into positive evidence. A pinned-checkpoint test reporting ok with the gate unset may simply have skipped, since a skip and a genuine pass both report ok under the default test harness. The same test reporting ok with MLXCEL_REQUIRE_PINNED_CHECKPOINTS=1 set can only mean the checkpoint was actually read and the contract asserted, because an unusable checkpoint would have failed instead. This is how acceptance criterion 2 of issue #1161 was proven, and it is a reason to reach for the variable even when nothing is suspected to be wrong.

Before PR #1173 introduced this gate, the availability pre-check was narrower rather than absent. PR #1157 had already replaced the vision-side test's assert!(index_path.exists(), ...) with a silent skip and added the same index-absent skip on the loading side, so a checkpoint with no index went quiet at that point. Everything past that guard still unwrapped: with the index present but config.json or a referenced shard missing or truncated, the test panicked on an unwrapped read error, which was loud and impossible to miss. PR #1173 replaced those unwraps with a pre-check that skips with a reason naming the offending file, while keeping genuine contract violations as failures, so a partially materialized checkpoint now reports ok quietly by default too. MLXCEL_REQUIRE_PINNED_CHECKPOINTS=1 is what buys the loudness back, for both the index-absent and the partially materialized case, on the one machine where the checkpoint is supposed to be complete.