Skip to content

MLXCEL Environment Variables

This page is the developer-facing inventory of MLXCEL_* shell environment variables used by mlxcel. It includes stable operator settings, service defaults, build controls, and low-level diagnostic switches.

For routine deployment, prefer documented CLI flags when a flag exists. Environment variables are process-wide, and several low-level switches are read once and cached on first use. Set them before starting mlxcel or mlxcel-server.

Conventions

  • CLI flags take precedence over matching environment variables unless the flag help says otherwise.
  • Some llama.cpp-compatible settings use LLAMA_ARG_* aliases. This page focuses on mlxcel-native MLXCEL_* names.
  • Boolean parsing is not uniform across every internal knob:
  • user-facing 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;
  • rows marked "falsy disables" treat 0, false, off, or no as disabled.
  • Variables marked advanced, diagnostic, or test-only are not stable public API. They exist for benchmark, rollback, kernel-development, or CI work.

Runtime and Sizing

Variable Values Default Notes
MLXCEL_DEVICE gpu, metal, cpu GPU hint Runtime device override. cpu requests CPU execution. Invalid values warn and fall back to GPU preference; if no GPU backend is available, the runtime falls back to the 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.
MLXCEL_WIRED_LIMIT max, 0, none, bytes, NGB, NMB max on Apple GPU Apple Silicon GPU wired-memory limit. 0 / none disables the limit; explicit sizes cap wired memory.
MLXCEL_WIRE_WEIGHTS 1/true/on/yes unset (off) macOS only: mlock the evaluated weight buffers after load, in 1 GiB chunks, stopping before the system wired total would pass RAM minus max(7 GiB, 15% of RAM). Needs the default eager MLXCEL_WEIGHT_MATERIALIZATION. Reports the wired total as mlxcel_weights_wired_bytes on /metrics and logs every failure. Off by default because no soak has run; see macOS Memory Guardrails.
MLXCEL_MEMORY_LIMIT 0, none, bytes, NGB, NMB unset Soft MLX allocator cap. Nonzero values also become the authoritative available-memory budget for mlxcel inspect and --estimate-memory.
MLXCEL_HEADROOM_FACTOR positive f64 1.20 Allocator-overhead multiplier on weights + KV in the unified memory estimator (mlxcel inspect, --estimate-memory, the paged KV auto budget, and the prompt-cache capacity defaults). The activation reserve (MLXCEL_ACTIVATION_MULT) and, on ROCm builds, the backend in-flight reserve (MLX_ROCM_MAX_INFLIGHT_MB) are separate terms that this factor does not scale. Positive values <= 1.0 disable the headroom term; invalid or non-positive values warn and fall back to the default.
MLXCEL_ACTIVATION_MULT positive finite f64 2.0 Multiplier on the activation working set (batch x min(context, prefill chunk) x (hidden + intermediate) x 2 bytes) in the unified memory estimator, with the same callers as MLXCEL_HEADROOM_FACTOR. MLXCEL_HEADROOM_FACTOR does not scale this term. This is a calibration override; invalid, non-finite, or non-positive values warn and use 2.0.
MLXCEL_CACHE_DIR directory path $HOME/.cache/mlxcel Root for the default model store ($MLXCEL_CACHE_DIR/models/<owner>/<name>) and other mlxcel disk caches, including tokenizer language-analysis files.
MLXCEL_MODELS_DIR directory path unset Direct model-store root override. Repo-id downloads and lookups use $MLXCEL_MODELS_DIR/<owner>/<name> with no extra models/ component. Overridden by --models-dir.
MLXCEL_DEFAULT_ORG HuggingFace owner mlx-community Default owner used when a model argument has no slash, for example Qwen3-4B-4bit -> mlx-community/Qwen3-4B-4bit.
MLXCEL_SERVER_DECODE_STORAGE auto, dense, paged auto Continuous-batching decode storage backend. The --decode-storage-backend flag takes precedence.
MLXCEL_SURGERY YAML file path unset Default runtime model surgery file when --surgery is not supplied and the surgery feature is built.
MLXCEL_CACHE_LIMIT 0, none, bytes, NGB, NMB 2GB on ROCm builds, unset elsewhere Bound on MLX's own buffer cache (set_cache_limit). Unset applies the backend default: no bound on Metal and CUDA, 2 GiB on ROCm, where the allocator's own default lets the cache grow to the allocator's memory limit. 0 or none removes the bound and leaves MLX's default cache behavior. On CUDA this is the main lever for bounding memory growth, since the periodic decode-loop cache clear defaults to off there; see MLXCEL_CACHE_CLEAR_INTERVAL below.
MLXCEL_CACHE_CLEAR_INTERVAL 0 (disable), positive integer (token cadence) 0 on CUDA, 256 on Metal/CPU Cadence, in generated tokens, of the periodic clear_memory_cache call in the decode loops and batch scheduler. The default is backend-aware: on CUDA the periodic clear churns the memory pool and defeats CUDA-graph reuse, so it is off by default there (bound the cache with MLXCEL_CACHE_LIMIT instead); on Metal/CPU the cheap 256-token trim mirrors Python mlx-lm.
MLXCEL_PREFILL_CHUNK unsigned integer tokens, 0 disables 2048 Prefill chunk size for every path (CLI, server, benches): the prompt is fed through the model in chunks of this many tokens, so sliding-window caches trim between chunks and prefill peak memory tracks the chunk size rather than the whole prompt. 0 prefills the whole prompt in one pass. On the server, setting it pins the chunk size when --prefill-chunk-size is not given, replacing the adaptive default (2048 while no other sequence is decoding, 512 while a decode batch is live) with this one value in both states; 0 disables chunking there too. An unparseable value logs a warning and keeps the default.
MLXCEL_MAX_BATCH_PREFILL_TOKENS unsigned integer tokens, 0 disables derived (2 * max_batch_prefill * min(prefill_chunk_size, 512), 4096 at the defaults) Padded-token budget bounding one server batched prefill window, which caps the transient memory a prefill batch can claim. The 512 cap keeps the 2048-token chunk default from quadrupling the budget. 0 removes the cap. --max-batch-prefill-tokens takes precedence.
MLXCEL_PREFILL_GRANT_INTERVAL unsigned integer decode ticks, 0 disables 16 Prefill fairness grant for the continuous-batching scheduler. A prompt longer than --prefill-chunk-size parks between chunks; every N decode ticks it is granted a chunk so the batch keeps decoding instead of stalling behind one long prompt. 0 removes the grant.
MLXCEL_KV_ADMISSION_WATERMARK fraction 0.0 to 0.5 0.01 Share of the paged KV block budget a prefill admission keeps free while rows are decoding, so they can grow without preemption. To make that room, admission evicts cold prompt-cache prefixes and preempts strictly lower-priority rows before it waits. Applies only while a decode batch is live and only with a block budget; 0 disables it. Out-of-range or non-numeric values are rejected at startup. --kv-admission-watermark takes precedence. See Continuous Batching.
MLXCEL_REQUEST_MEMORY_BUDGET positive bytes, NK/NKB, NM/NMB, NG/NGB smallest positive runtime cap (device recommended working set, wired cap, MLX allocator cap) Request admission ceiling for text-only Nemotron-H on Metal, where the scheduler serves one active slot and admits each request against a conservative memory estimate. The variable can only lower the runtime ceiling. Explicit zero, none, empty, nonfinite, or overflowing values stop server startup. An over-budget request gets HTTP 503 request_memory_budget_exceeded before prefill (and before SSE headers on a streaming request).
MLXCEL_PREFILL_DEQUANT_MIN_M 0/off/false/no disables; positive integer row count 1024 on M1-generation Apple Silicon, off on other Apple generations, CUDA, and ROCm Input row count from which an affine 4-bit projection whose scales share the input dtype (f16 or bf16) runs as dequantize plus a dense matmul instead of quantized_matmul. Only projections with a wide enough output qualify, where both routes return identical bytes. On M1 Ultra this raises prefill at 1024 rows by up to about 15% depending on the model, at a cost of 0.1 to 0.6 GB of prefill peak memory.
MLXCEL_MIXED_STEP 1/true/yes/on to enable; unset or anything else to disable off Experimental. Mixed prefill/decode step prototype for the continuous-batching scheduler, letting one tick carry both work types. See ADR 0005 for the design and its rejected alternatives.
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 kernel, the fallback splits its score matrix to stay inside this budget.
MLXCEL_ALLOWED_ORIGINS comma-separated origin list unset (permissive) Restricts CORS to the listed origins, for example https://app.example.com,https://admin.example.com. The --allowed-origins flag takes precedence.
MLXCEL_WEIGHT_MATERIALIZATION eager, lazy (case-insensitive) eager When a generation model's checkpoint tensors are read and uploaded. eager evaluates them at the end of model load, so Model loaded in ... reports resident weights and the first token pays no weight I/O; a server pays it at startup instead of on its first request. Modality-tower tensors (vision_tower, audio_tower, multi_modal_projector and similar) stay lazy and load on the first image or audio request. lazy materializes every weight inside the first evaluation that reaches it. Tensor-parallel and pipeline-parallel loads and embedding or reranker loads always materialize on first use. Invalid values warn and use eager.
MLXCEL_DENSE_GEMV_ROWS integer row cap; 0 disables 6 On Apple GPU generations 13 and 14 (M1, M2), dense matmuls with up to this many rows run the GEMV kernel as a batch instead of the steel GEMM tile. 0 restores upstream dispatch. Read once per process.

Related operator docs: Configuration, CLI Reference, Memory Estimation, and Runtime Model Surgery.

MLX Runtime Variables

These non-MLXCEL_* values are initialized by the CLI unless the operator set them explicitly.

Variable Values Automatic default Notes
MLX_MAX_OPS_PER_BUFFER positive integer 1000 on M1-M4; 100 on CUDA compute capability 12.1 (GB10) for laguna checkpoints; 2000 on ROCm (the backend's own default); MLX default elsewhere Ops captured into one Metal command buffer, one CUDA graph, or one ROCm command batch before it is committed. On CUDA it binds together with MLX_MAX_MB_PER_BUFFER, since MLX commits when either budget is exceeded; on ROCm a batch is also committed once it has allocated a quarter of MLX_ROCM_MAX_INFLIGHT_MB. The GB10 default is a family allowlist because raising both budgets helped Laguna on every measured workload and slowed other families; it also raises peak memory. An explicit value always wins.
MLX_MAX_MB_PER_BUFFER positive integer (an element count shifted right by 20, not bytes) Metal: MLX default, raised for decode steps only by MLXCEL_DECODE_MB_PER_BUFFER; CUDA: 1000 on compute capability 12.1 for laguna checkpoints; MLX default elsewhere Input-size budget of one command buffer or CUDA graph. Setting it pins the value for prefill and decode alike and disables the decode-only raise. An explicit value always wins.
MLX_CUDA_SDPA_CACHE_SIZE unsigned integer 2000 on CUDA (MLX's own default is 256) LRU capacity of MLX's cuDNN SDPA execution-plan cache. Each distinct attention shape is a miss, and MLX aborts with Cache thrashing once lifetime misses pass twice the capacity, so a long-lived server with diverse prompt lengths needs the larger budget. An LRU cap, not a preallocation. An explicit value always wins.
MLX_CUDA_GRAPH_CACHE_SIZE positive integer 2000 on CUDA Raises MLX's graph-cache LRU capacity for long-lived shape-diverse decode. An explicit value always wins.
MLX_USE_CUDA_GRAPHS 0 disables; other MLX-supported values enable MLX default; auto-0 for Gemma 4 on CUDA Controls MLX CUDA graph capture. mlxcel disables it before first evaluation for Gemma 4 to avoid the known incremental-decode hazard; an operator-supplied value in either direction always wins.
MLX_ENABLE_TF32 0 (full f32), 1 (TF32-class) 1 (MLX's default; mlxcel does not override it) MLX's own switch for the precision of f32 GEMMs on CUDA and Metal. With the default, f32 GEMM inputs go through a 10-bit mantissa: cuBLAS TF32 compute on CUDA Ampere and later, the NAX kernel on Apple GPU generation 17. It does not affect bf16, f16, or quantized matmuls. Set 0 when an f32 checkpoint needs f32-exact output; the repository's test binaries pin it to 0.
MLX_PTX_CACHE_DIR directory path $MLXCEL_CACHE_DIR/cuda-ptx/<MLX-pin> on CUDA Persistent NVRTC PTX cache. mlxcel creates a directory scoped to the pinned MLX commit so stale kernels do not survive an MLX upgrade. An explicit path always wins.
MLX_BF16_NATIVE 0 selects legacy conversion; anything else keeps native bf16 native bf16 Low-level bridge compatibility switch. 0 expands bf16 tensors to f32 at load for hardware/configurations that cannot use the native bf16 path, increasing memory use.
MLXCEL_CCCL_DIR existing directory path bundled/build-tree headers Runtime NVRTC include override for CCCL headers. Intended for relocatable CUDA packages and embedders.
MLXCEL_CUTLASS_DIR existing directory path bundled/build-tree headers Runtime NVRTC include override for CUTLASS/CuTe headers.
MLXCEL_QUIET_JIT presence suppresses unset Suppresses the one-time cold-cache CUDA kernel compilation notice; it does not disable JIT compilation.

External Model Storage

These control where a Qwen3.8-Flash-Next (qwen4_exp) checkpoint keeps its n-gram table and routed experts. See External Weight Stores.

Variable Values Default Notes
MLXCEL_PLE_IO direct, mmap, ram mmap I/O path for a checkpoint whose n-gram table is external (text_config.ple_storage); any other value is a load error. mmap maps the row file and copies the selected rows through the page cache. direct reads rows unbuffered through a small thread pool (F_NOCACHE on macOS, O_DIRECT where available on Linux) and coalesces neighbouring rows. ram reads the whole table into anonymous memory once and keeps no row cache. Ignored without ple_storage; an absolute manifest path, or one that resolves outside the model directory, is refused at load.
MLXCEL_PLE_CACHE_BYTES unsigned bytes, 0 disables 536870912 (0.5 GiB) Row-cache budget for an external n-gram table. When unset, a positive cache_rows in ple_storage or the manifest sizes the cache as cache_rows times the row stride. Unused in ram mode.
MLXCEL_EXPERT_STREAM 1/on/true, 0/off/false, auto auto SSD expert streaming for a Qwen3.8-Flash-Next checkpoint on Metal: routed experts live in a resident slot arena and are read from the SSD as tokens route to them. auto engages only when the routed experts exceed 60% of Metal's working set, and falls back to the resident path for a layout it cannot stream; 1 forces it and makes such a layout a load error. Output stays exact. Other MoE families, CUDA, and ROCm are not covered.
MLXCEL_EXPERT_ARENA_BYTES unsigned bytes derived from free memory Arena size for MLXCEL_EXPERT_STREAM, split across the layers at 64 to num_experts slots each. Unset, the arena is the memory macOS can give now minus the resident non-expert weights and a 6 GB headroom, capped at the experts' size.

Downloader

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_MODELS_DIR directory path unset Store root used by download, repo-id auto-resolution in run/generate/serve/inspect, list, and rm. CLI --models-dir wins.
MLXCEL_DEFAULT_ORG HuggingFace owner mlx-community Owner prepended to bare model names.
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 custom CA certificates on top of the operating system trust store that the download client verifies against. The server's remote-media fetcher uses the same client configuration. Point it at a PEM bundle for hosts behind a TLS-inspecting proxy or a private registry.

Prompt Cache

These variables apply when the corresponding CLI flag is absent. See Prompt Cache for behavior and tuning.

Variable Values Default Flag equivalent
MLXCEL_PROMPT_CACHE_ENABLED boolean true --prompt-cache-enabled
MLXCEL_PROMPT_CACHE_CAPACITY_BYTES unsigned bytes model-aware, never below 2147483648 --prompt-cache-capacity-bytes
MLXCEL_PROMPT_CACHE_MAX_ENTRIES unsigned integer 1024 --prompt-cache-max-entries
MLXCEL_PROMPT_CACHE_TTL unsigned seconds 3600 --prompt-cache-ttl
MLXCEL_PROMPT_CACHE_MIN_PREFIX unsigned tokens 32 --prompt-cache-min-prefix
MLXCEL_PROMPT_CACHE_SNAPSHOT_CAPACITY_BYTES unsigned bytes model-aware, never below 536870912 --prompt-cache-snapshot-capacity-bytes. Byte budget for the session snapshot store, accounted separately from the prefix cache.
MLXCEL_PROMPT_CACHE_SNAPSHOT_MAX_ENTRIES unsigned integer 4096 --prompt-cache-snapshot-max-entries
MLXCEL_PROMPT_CACHE_SNAPSHOT_TTL unsigned seconds 7200 --prompt-cache-snapshot-ttl
APC_ENABLED boolean true --apc-enabled; APC is automatically disabled for incompatible recurrent/hybrid models
APC_BLOCK_SIZE unsigned tokens 16 --apc-block-size
APC_NUM_BLOCKS unsigned integer derived --apc-num-blocks
APC_HASH sha256, blake3 sha256 --apc-hash

MLXCEL_PROMPT_CACHE_ENABLED wins over the llama.cpp-compatible LLAMA_ARG_CACHE_REUSE alias when both are set and no CLI flag is provided.

When neither the flag nor the variable sets a capacity, both stores are sized from the loaded model. The snapshot store gets max(536870912, min(6 * entry, left / 4)) for snapshot-only families (recurrent hybrids such as Qwen 3.5, Jamba, Falcon-H1, LFM2, Nemotron-H, Mamba, Granite 4 hybrid, PLaMo 2, and attention-cache families such as Gemma 3, Gemma 4, Llama 4, AFMoE), where entry is one entry at a representative length (--ctx-size, or 8192) and left is the memory left after the model's own estimated footprint; it is raised to one whole entry when that fits in left / 2. The KV store gets the same treatment for dense-KV families and never drops below the 2 GiB fallback. An entry larger than its whole store logs one warning per store that names the flag to raise.

Variable Values Default Flag equivalent
MLXCEL_ENABLE_VLM_PREFIX_CACHE boolean false --enable-vlm-prefix-cache. Experimental VLM prompt-prefix cache sharing for multi-turn same-image conversations.
MLXCEL_DISABLE_BOUNDARY_SNAPSHOT exactly 1 disables; anything else, including true and on, leaves it enabled on Diagnostic kill switch. Stops prefill from donating a snapshot at the history boundary. That donation is what makes multi-turn prompt caching hit at all for snapshot-only families (Gemma 4, Qwen 3.5, the SSM hybrids), so it is on by default.
MLXCEL_DISABLE_CACHE_WARMUP exactly 1 disables; anything else leaves it enabled on Diagnostic kill switch. Stops the background warm-up that 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.
MLXCEL_APC_TRACE exactly 1 enables; anything else, including true and on, leaves it off off Diagnostic. One trace line per scheduler prompt-cache decision, including which snapshot a request matched and why a divergent one was rejected. Read once at first use, so setting it after the process starts has no effect.

Speculative Decoding

Variable Values Default Notes
MLXCEL_DRAFT_KIND dflash, mtp, full-model; on mlxcel serve and mlxcel-server also prompt-lookup auto / none Alias for --draft-kind when the CLI flag and LLAMA_ARG_DRAFT_KIND are absent. full-model drafts with an ordinary smaller checkpoint named by --draft-model / --model-draft. mlxcel generate takes prompt lookup through --prompt-lookup instead and rejects prompt-lookup here.
MLXCEL_DRAFT_BLOCK_SIZE unsigned integer resolved at startup Alias for --draft-block-size when the CLI flag and LLAMA_ARG_DRAFT_BLOCK_SIZE are absent. When all three are unset, the width is a width the drafter declares, then a measured per-device default (DFlash with an affine-quantized target on GB10 resolves to 4), then 4 for MTP and full-model and 16 for DFlash. The startup speculative= line logs the width and the rule that produced it. Worker-owned: a change needs a restart.
MLXCEL_DFLASH_PER_ROW_ACCEPTS 0/false/no/off to disable; unset or any other value to enable on Per-row accept commits in the batched Qwen 3.5 DFlash round loop. Each row commits its own accept count instead of clamping the round to the shortest live row: the attention caches compact back to a per-row contiguous prefix, queries and keys rotate at each row's own positions with its stale gap masked, and each linear-attention layer replays its convolution and recurrent state over the accepted prefix. The clamp returns with 0 for A/B work, and is also taken automatically when the target's attention caches are quantized, paged, or head-trimmed, because those cannot be compacted row by row.
MLXCEL_ENABLE_MTP_B1 truthy / falsy adaptive Pins the B=1 MTP decision and skips adaptive profiling. 0, false, no, or off force classic decode; other values force MTP. When unset, MLXCEL_MTP_ADAPTIVE profiles each pairing and uses the older per-hardware rule only as an ambiguous-result fallback.
MLXCEL_ENABLE_MTP_BATCH truthy value off Advanced. Forces the batched Gemma 4 MTP burst path for parity/debug testing. On row-wise geometries (Gemma 4 31B, and the 12B on CUDA) the batched burst is decode-exact and runs on CUDA for windows of up to 4 rows whose prompts fit the sliding window; other windows, and Metal and ROCm, decline to classic decode. A window forms only from requests queued together that share a seed and carry no token bias, and only with --max-batch-prefill 1.
MLXCEL_ENABLE_MTP_BATCH_RAGGED truthy value off Advanced. Allows different prompt lengths in a batched MTP window. Requires MLXCEL_ENABLE_MTP_BATCH; declines prompts beyond the target sliding window.
MLXCEL_MTP_DISABLE_DIVERGENT_FIX 1 off Advanced. Rolls back the per-row compaction applied after divergent accepts in batched Gemma 4 MTP. Diagnostic only.
MLXCEL_MTP_BLOCK_CONTROLLER proxy, requested, or unset for the throughput comparator throughput Selects which controller decides the B=1 MTP verify width when --draft-block-size exceeds the drafter's configured depth. The default alternates 32-round measurement windows between the configured depth and the requested ceiling and holds whichever emits more tokens per millisecond, requiring a 2% lead to displace the incumbent and re-challenging the loser on a growing backoff. Evidence lives in the generator, so a server process keeps it across requests. proxy restores the fully-accepted-prefix gate, which holds the Gemma 4 12B pairing about 5% below its optimum because at 0.585 acceptance the configured prefix is rarely fully accepted. requested drafts at the requested width from the first round with the harmful-ceiling protection off, which is a width-sweep measurement mode rather than a deployment setting. Drafters that set prefer_requested_block_size (Qwen 3.5 MTP) bypass both controllers, and the batched B>1 loop stays on the proxy gate. Read once per generator.
MLXCEL_MTP_ADAPTIVE 0/false/no/off to disable, any other value or unset to enable on Adaptive B=1 MTP policy. The server profiles four qualifying B=1 bursts per pairing, persists an enable/decline verdict, and uses the static per-hardware table only when the measurement is ambiguous.
MLXCEL_MTP_TICK_SLICE 0/false/no/off to disable, any other value or unset to enable on Tick-cooperative B=1 MTP serving. A B=1 MTP request yields the scheduler between slices, so a speculative burst does not block the head of the queue.
MLXCEL_MTP_SLICE_GRANT_ROUNDS non-negative integer 8 Grant budget for one hold of the tick-slice speculative slot, counted in executed slices. When the budget is spent the slot rotates to another waiting request. 0 disables rotation and lets the active request keep the slot.
MLXCEL_MTP_QUANTIZE_DRAFTER 0/false/no/off to disable, any other value or unset to enable on Quantizes the Qwen 3.5 MTP drafter's projections at load, taking a bf16 drafter from about 810 MiB to about 228 MiB. The target verifies every proposal, so this trades acceptance rather than correctness, and acceptance did not move measurably. Read once per process.
MLXCEL_MTP_ALLOW_INEXACT 1/true/yes/on to enable; unset or anything else to disable off Engages Qwen 3.5 or Gemma 4 MTP even when the startup exactness probe reports that the multi-token verify block is not byte-identical to the single-token chain. Accepting it means temperature-0 speculative output may differ from non-speculative decode. Read once per process.
MLXCEL_MTP_PROBE_PROMPT_LEN integer >= 2 8 Advanced. Prompt length the Qwen 3.5 block-versus-chain exactness probe prefills before comparing. Kept so the length sweep can be repeated on another checkpoint; longer values cost one prefill per probe draw at worker startup.
MLXCEL_GLM_MTP_HIDDEN_TAP pre (also pre_norm, pre-norm); unset or anything else for the default post-final-norm Which target hidden state feeds the GLM-4.7-Flash (glm4_moe_lite_mtp) MTP drafter. pre exists for acceptance measurements; it changes the acceptance rate, never the emitted tokens. Read once per process.
MLXCEL_PRINT_TOKEN_IDS presence enables unset Diagnostic. Offline mlxcel generate prints the prompt ids and the generated token ids to stderr as [prompt ids (N): ...] and [token ids (N): ...], so a plain run and a --draft-model run can be compared on ids.
MLXCEL_QMV_WIDE any value pins the selection unset Advanced. Pins MLX's qmv_wide kernel selection yourself. An explicit setting wins over the exactness gate's own retry, so the retry is skipped and a decline reports that the byte-identity contract was unreachable under your pin.
MLXCEL_GDN_CHAIN_PARITY 0 to disable, any other value or unset to enable on Advanced, diagnostic. Gates the Metal chain-parity gated-delta kernel used by Qwen 3.5 MTP's verify and rollback-replay paths. Setting it to 0 forfeits the byte-identity contract; it exists for A/B attribution of the parity kernel's own cost. Metal-only.
MLXCEL_MTP_TREE_BRANCHES non-negative integer 0 Advanced. Runner-up candidates a tree round may spend verify positions on. 0 proposes the linear tree, which is exactly the chain behavior and the only byte-identical one.
MLXCEL_MTP_TREE truthy value off Enables MTP tree-round construction when the target reports support. Off by default; the chain path remains the measured production baseline.
MLXCEL_MTP_TREE_MARGIN float no gate Advanced. Top-two logit gap under which a step counts as uncertain enough to spend a branch budget on.
MLXCEL_MTP_DRAFT_PROFILE 1/true/yes/on to enable; unset or anything else to disable off Diagnostic. Evaluates each MTP drafter-step component before the next begins, so the MTP drafter step profile log line reports real GPU time instead of graph-build time. The extra syncs inflate the step cost, so quote the unprofiled number as the cost and the profiled one only for attribution.
MLXCEL_MTP_GAP_LOG truthy value off Diagnostic. Prints position, token, top1, top2, gap at each verify position, for judging whether the exactness cost could be made conditional on near-ties.

See Speculative Decoding for the supported user-facing paths.

Loop Detection

Variable Values Default Notes
MLXCEL_LOOP_DETECTION off/0/none/false/disabled, on/default/true/enabled, or MIN,MAX,COUNT unset (per-family auto-enable applies) Global override for tail N-gram repetition ("loop") detection. Unset leaves the per-family auto-enable policy in place (on by default for the Gemma 4 family). off force-disables for every request; on forces the recommended threshold (min=1, max=20, count=4); an explicit triple such as 1,20,4 sets min_pattern_size,max_pattern_size,min_count directly. A per-request max_pattern_size / min_pattern_size / min_count in the API body always wins over this variable. A malformed value warns and is ignored.

Loop detection ends generation early when a short token pattern repeats at the tail of the stream, the countermeasure vLLM exposes in SamplingParams. See REST API Reference for the per-request fields.

Block Diffusion

Variable Values Default Notes
MLXCEL_DIFFUSION_DEBUG_CANVAS 1 off Test-only. Replaces random canvas initialization with a fixed deterministic pattern for cross-implementation parity testing. Output is not useful for normal generation.

See Block Diffusion for the generation mechanism and the --diffusion-* CLI flags.

KV Cache and TurboQuant

Prefer CLI flags such as --cache-type-k, --cache-type-v, --kv-cache-mode, --turbo-boundary-v, and the server batch KV quantization flags when possible.

Variable Values Default Notes
MLXCEL_KV_CACHE_BUDGET auto, bytes, none, or 0 auto Block budget for the unified paged KV cache pool. auto derives a memory-aware cap; none/0 opts out. Alias for --kv-cache-budget; applies to pool-backed FP16 models under the paged decode storage backend.
MLXCEL_KV_BOUNDARY_V_LAYERS integer count 2 Number of first/last V layers kept at higher precision for Turbo4-family modes. 0 disables. --turbo-boundary-v writes this before cache construction.
MLXCEL_TURBO_BOUNDARY_V integer count alias Compatibility alias for MLXCEL_KV_BOUNDARY_V_LAYERS; the primary name wins when both are set.
MLXCEL_KV_SKIP_LAST_LAYER boolean true mlxcel-native 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_COUNT output file path unset (off) Diagnostic. When set to a non-empty path, every single-token Turbo4Asym decode step appends a call_idx,kv_tokens,skipped,total row to that CSV, counting post-softmax attention weights below MLXCEL_SPARSE_V_THRESHOLD across all heads. A graph-only side computation; zero cost when unset. See the measurement assets in benchmarks/sparse_v_skip/.
MLXCEL_SPARSE_V_KERNEL falsy disables enabled on macOS Allows fused Sparse-V/dequant Metal kernels. Set 0, false, off, or no to force graph fallback.
MLXCEL_TURBO4_DEQUANT_SDPA falsy disables on Controls the dequant-first SDPA path for symmetric Turbo4.
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_KV_INPLACE_WRITE 0/false/off/no disables on A single-token FP16 decode step writes its K/V row into the dense cache buffer in place instead of through slice_update, which otherwise copies the whole cache every step while the previous pipelined step still holds the buffer. The in-place write is taken only when the cache's own last write produced the current buffers, so a buffer shared with a prompt-cache snapshot is never rewritten. The server's paged pool takes the same in-place write for rows of a block owned by one sequence. Measured on M1 Ultra: command-r7b 4-bit decode +2.5% at context 16 to +8.2% at 2048.
MLXCEL_DIAG_SKIP_DECODE_KV_WRITE presence enables off Diagnostic; output is invalid while set. A single-token FP16 decode step skips writing its K/V row, so a decode benchmark reads the upper bound of removing the per-step cache write. The cache never grows and positions do not advance, so the generated text is meaningless. Prints one warning when first used.
MLXCEL_TURBO4_ASYM_DEQUANT_SDPA falsy disables on Dequant-first SDPA path for asymmetric Turbo4Asym (FP16 K + 4-bit V). Falsy values select the lossy sparse-V path instead.
MLXCEL_MLA_ABSORBED 1/true/on/yes enables off Matrix-absorbed MLA decode for the DeepSeek family. Caches the compressed latent (ckv, kpe) rather than decompressed per-head K/V, which cuts KV footprint substantially on long contexts.
MLXCEL_MLA_SPLIT_KV 1/true/on/yes enables off Advanced. Splits the latent range into chunks whose partial softmax states are merged by the split-KV merge kernel. Requires MLXCEL_MLA_ABSORBED.
MLXCEL_CASCADE_ATTENTION 1/true/on/yes enable; 0/false/off/no kill switch off Two-level cascade decode. A whole-page prompt prefix shared by several concurrent sequences is attended once, then merged with each request's private suffix state, removing the duplicated work of a shared system prompt.
MLXCEL_CASCADE_MIN_SHARED_PAGES non-negative integer 16 Whole pages a subgroup must share before cascade is used. At the default block size of 32, 16 pages is 512 tokens. Below the threshold the flat launch is used.
MLXCEL_CASCADE_MIN_MEMBERS integer >= 2 2 Sequences that must share the span before cascade is used. Below two there is no duplicated work to remove.
MLXCEL_NATIVE_WINDOW_SDPA auto, 1/on, 0/off auto Advanced, A/B selector. Chooses how a causal multi-query prefill over a sliding window reaches MLX's SDPA on Metal once the returned keys outgrow the window: MLX's native window_size argument with mask_mode = "causal", or the explicit full-width windowed-causal mask the array-mask path always took. auto takes the native route only where MLX reports a fused kernel for the shape; 1 forces the native entry even through the graph fallback; 0 keeps the legacy mask route. Only statically eligible Metal prefill shapes are affected (no softcap, no attention sink, no extra mask array, query block longer than 8 rows and no longer than the keys, supported head dims, f16/bf16/f32); decode, small verify blocks, CPU, CUDA, and ROCm keep their existing route. Read once per process.
MLXCEL_LOG_WINDOW_SDPA 1/sampled, all off Prints one [mlxcel][window-sdpa] line per causal window prefill to stderr with the route taken (native or legacy), the reason, MLX's fused-kernel verdict for both arms, and the q/k shapes. sampled (the same as 1) logs the first 16 calls and then every 256th. Read once per process.

Video and Local Media

These apply to video-capable VLM request handling. See Vision-Language Models.

Variable Values Default Notes
MLXCEL_ALLOW_PRIVATE_MEDIA_URLS 1 to enable unset (refusal active) Turns every network-address check on the server's remote media fetches into a pass, for a deployment that genuinely serves its 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 Server.
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 the directories owner-writable only; group- or world-writable entries warn at startup. Both native-video families and the video-frames fallback resolve the list once when the server starts (once per model the router loads, and once for an in-process mlxcel run), so an entry that does not exist yet is dropped with a warning. Restart after changing the variable or creating a listed directory. Data-URI and HTTP(S) video payloads are written as create-new 0600 files inside private 0700 directories before ffmpeg reads them.
MLXCEL_VIDEO_MAX_FRAMES unsigned integer (minimum 2) 16 Frames kept when a clip is sent as ordered still images to a checkpoint with no native video path. Values below 2 are raised to 2. Same as --video-max-frames on mlxcel serve, mlxcel-server, and mlxcel generate. The frames count against --max-images.
MLXCEL_VIDEO_FPS float > 0 2.0 Sampling rate of that fallback when the request's video_url.fps is absent. Server only (--video-fps); the CLI uses --fps.
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 bytes 268435456 Per-frame cap for the ffmpeg PNG stream splitter.
MLXCEL_MEDIA_NAT64_PREFIXES comma-separated IPv6 prefixes addr/len, with len one of 32, 40, 48, 56, 64, 96 unset (no prefixes) NAT64 prefixes the media address policy decodes. An address inside a listed prefix is judged by the IPv4 address it embeds (RFC 6052), so a synthesized public target is fetched while a synthesized private one stays refused. The well-known 64:ff9b::/96 is always decoded. A listed prefix must lie inside global unicast 2000::/3 or the RFC 8215 local-use block 64:ff9b:1::/48; an entry that does not parse, uses another length, or falls outside those ranges fails startup naming this variable. Read once at startup; MLXCEL_ALLOW_PRIVATE_MEDIA_URLS=1 still disables every address check.

MLXCEL_KIMI_K3_MAX_MEDIA_TOKENS (default 32768) caps the media tokens one Kimi K3 request's images may use together, because a single image can reach nearly 17,000 tokens under the checkpoint's resize rule and the tower's attention is quadratic in each image's patch count.

The server's image admission limits also have llama.cpp-compatible aliases: LLAMA_ARG_MAX_IMAGE_WIDTH and LLAMA_ARG_MAX_IMAGE_HEIGHT (both default to 16384). See the serve CLI reference.

Disaggregated Router Security

Variable Values Default Notes
MLXCEL_DECODE_ALLOWLIST comma-separated numeric IP:port entries unset (permissive with warning) Full set of router-selectable decode nodes that a prefill worker may connect to for KV handoff. Invalid/hostname entries are skipped. Configure it on a trusted cluster to prevent a forged decode_target from redirecting prompt-derived KV state.
MLXCEL_ROUTER_STATS_VERBOSE 1/true/yes/on enables off Includes raw peer addresses in GET /router/stats. The default redacts addresses and exposes stable router node ids instead.
MLXCEL_CLUSTER_SECRET_FILE path to a file holding the cluster secret unset Pre-shared per-cluster secret for distributed pipeline roles and disaggregated roles. The file must hold at least 16 bytes (one trailing newline is ignored); generate it with openssl rand -hex 32 > cluster.secret and copy it to every host. Precedence: --cluster-secret-file, then this variable, then MLXCEL_CLUSTER_SECRET, then [cluster] secret_file in the cluster TOML; setting both environment variables is an error. Every pipeline and disaggregated role refuses to start without a secret, and single-node serving does not read it.
MLXCEL_CLUSTER_SECRET the secret value (at least 16 bytes) unset Inline form of MLXCEL_CLUSTER_SECRET_FILE for containers and service units that inject secrets through the environment. Same precedence and fail-closed behavior, and the value is never logged. Prefer the file form on shared hosts because environment variables are visible to other processes of the same user; setting both environment variables is an error.

OpenXLA Backend

The root Cargo manifest declares xla-backend, xla-iree and the xla-diagnostics variants, and every one of them is default-off. Stock Metal and CUDA binaries therefore compile none of the OpenXLA path and do not read these names; a build that passes --features xla-backend (the seam and the StableHLO emitter, no native toolchain needed) or --features xla-iree (the seam plus a real IREE runtime, which needs IREE_DIST, IREE_CUDA_HOME, IREE_MACOS_HOME, or IREE_ROCM_HOME at build time) does. See OpenXLA Backend Source Status.

Variable Values Default Notes
MLXCEL_BACKEND xla MLX (unset) Selects the OpenXLA/StableHLO backend. Inert in a stock binary, which is built without xla-backend.
MLXCEL_XLA_DEVICE local-task, local-sync, metal, cuda, hip metal on Apple Silicon, local-task elsewhere IREE HAL device for the OpenXLA engine. local-task / local-sync are CPU drivers. cuda and hip are never auto-selected; they need a runtime built with that driver and must be set explicitly.
MLXCEL_XLA_LOCAL_TASK_WORKERS integer, 1 to 256 unset (one worker per physical core of the current NUMA node, at most 64) Worker count of the local-task (CPU) device. The shim builds that device itself with 8 MiB worker stacks, because IREE's 128 KiB default stack cannot hold the static TLS of a process that has loaded the CUDA or ROCm runtime libraries. Read only when the device is local-task; 0, a negative, a non-integer, or a value above 256 fails the load with an error naming the variable instead of being clamped.
MLXCEL_XLA_ROCM_TARGET one gfx* chip, for example gfx1151, gfx1100, gfx942 unset Required with MLXCEL_XLA_DEVICE=hip; passed to iree-compile as --iree-rocm-target. Unset, empty, comma-separated, or non-gfx* values are rejected before iree-compile runs. scripts/iree/setup-rocm.sh --env exports the detected host target.
MLXCEL_XLA_PRECISION f32, f16, bf16 f16 on GPU devices (metal, cuda, hip), f32 on CPU devices Contraction precision for the StableHLO emitter's matmuls; norms, softmax, and RoPE stay f32 regardless. bf16 is rejected at load on a metal device (no bf16 compute on Metal GPU targets). An unrecognized value is rejected with a [compile] diagnostic.
MLXCEL_XLA_QUANT packed unset Keeps MLX 4/8-bit quantized weights packed inside the graph with device-side dequantization, instead of dequantizing at load. Supported on cuda, hip, and local-task; rejected at load on a metal device, where the packed prefill faults. Token-exact but off by default: IREE lowers the in-graph dequant and the matmul as separate dispatches, so host-dequantized weights decode faster, while packed cuts host upload bytes.
MLXCEL_XLA_VISION_BACKEND auto, iree, host auto Vision tower and projector selector for the shared CLI/server prepared-prefill path. auto uses the resident IREE module when one is available and falls back to the host (MLX) path otherwise; iree fails startup instead of falling back; host keeps vision on MLX. Qwen2-VL and Qwen2.5-VL have no host vision path, so they require an xla-iree build and reject host.
MLXCEL_XLA_CONTEXT_CAPACITY integer tokens in 1..=2147483647 unset (request-sized on the CLI, derived buckets on the server) When set, pins exactly one static sequence capacity compiled into the OpenXLA prefill/decode pair and every KV buffer, and rejects a request before native execution when effective_prompt_len + max_new_tokens exceeds it. When unset, the CLI sizes the shape from the request: the smallest rung of [256, 512, 1024, 2048, 4096, 8192, 16384, 32768] that admits the prompt plus the generation budget, capped by the checkpoint's max_position_embeddings; an unlimited -n -1 sizes with a 512-token budget and clamps the admitted budget to the room the shape leaves. A larger capacity grows KV memory linearly and the prefill attention mask quadratically, and compiles a distinct artifact. StableHLO dynamic sequence shapes are not used.
MLXCEL_XLA_CONTEXT_BUCKETS comma-separated integers, each in 1..=2147483647 derived (256, min(4096, checkpoint window), plus the worst-case image floor when larger) Server OpenXLA continuous-batching worker only (xla-iree): the static context shapes it compiles, sorted and deduplicated, for an operator who knows the traffic better than the derivation. Requests route to the smallest bucket that admits them, so short requests keep the 256 shape. Ignored when MLXCEL_XLA_CONTEXT_CAPACITY is set, which yields a single bucket. An empty list, a 0, or a non-integer aborts worker startup. The CLI path does not read it.
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 and ROCm source-runtime builds bake the matching compiler from IREE_CUDA_COMPILE or IREE_ROCM_COMPILE.
MLXCEL_XLA_LOAD_TIMING presence enables off Diagnostic. Prints mlxcel-xla-load: phase=<label> elapsed_ms=<n> rss_kib=<n> lines to stderr for each phase of the OpenXLA language-model load and the Gemma3n audio runtime load. RSS comes from /proc/self/status, so it reads unavailable off Linux.
MLXCEL_PHI4MM_AUDIO_PRECISION f32, bf16 f32 Phi-4 multimodal audio encoder precision in the retained XLA source. Invalid values fail graph preparation.
MLXCEL_XLA_GEMMA3N_TRACE_LAYER_START non-negative layer index unset Diagnostic start layer for Gemma 3n host/XLA trace comparison.
MLXCEL_XLA_REFERENCE_EXPECT_UNEXPANDED_IDS test/reference policy value unset Internal XLA reference-admission contract switch; not a stock-server setting.
MLXCEL_XLA_IREE_RUN_MODULE executable path discovered from IREE tooling Test/tool override for iree-run-module in retained Gemma 3n probes.
MLXCEL_XLA_IREE_TEST_TARGET local or probe target local Target selector for ignored retained-backend integration probes.

Server Audio Admission

Variable Values Default Notes
MLXCEL_AUDIO_QUEUE_DEPTH unsigned integer 8 Bound on the audio worker command queue. When the queue is full, new audio requests receive a structured 503. The --audio-queue-depth flag takes precedence.
MLXCEL_AUDIO_REQUEST_TIMEOUT_SECS unsigned integer seconds 120 Per-request reply timeout, so a stuck or pathologically slow audio request frees its queue slot instead of holding it. The --audio-request-timeout-secs flag takes precedence.

Server Stores and Request Limits

Variable Values Default Notes
MLXCEL_DECODE_TIMEOUT unsigned integer seconds 600 Per-request decode watchdog: cancels a request whose generation stops producing tokens. The --decode-timeout flag takes precedence. --timeout / LLAMA_ARG_TIMEOUT is a different control: the HTTP socket read/write timeout with llama-server's 3600-second default. Setting --timeout without also setting the decode watchdog logs a migration warning at startup. 0 falls back to the built-in 300-second guard with a warning rather than expiring every request.
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.
MLXCEL_CONVERSATION_STORE_MAX_BYTES unsigned integer bytes 67108864 Bounds the approximate retained JSON bytes of the conversation store. The --conversation-store-max-bytes flag takes precedence.
MLXCEL_RESPONSES_STORE_MAX_BYTES unsigned integer bytes 268435456 Bounds the approximate retained JSON bytes of the Responses API store. The --responses-store-max-bytes flag takes precedence.
MLXCEL_RESPONSES_BACKGROUND_MAX_JOBS unsigned integer 64 Cap on concurrently active background: true Responses API jobs, each holding one generation thread. 0 disables background mode. The --responses-background-max-jobs flag takes precedence.
MLXCEL_RESPONSES_BACKGROUND_EPHEMERAL_TTL_SECS unsigned integer seconds, at least 1 600 How long the final object of a store: false background job stays readable by polling. 0 fails startup. The --responses-background-ephemeral-ttl-secs flag takes precedence.
MLXCEL_ANTHROPIC_PIN_BILLING_STAMP 0 disables (exact match only); any other value or unset leaves pinning on unset (pinning on) Pins the volatile cch and cc_version parts of Claude Code's x-anthropic-billing-header in the Anthropic Messages system text, replacing them with f characters of the same length before rendering, so a turn reuses the previous prompt's cached prefix. Only top-level system text starting with the exact header is touched. Read once through a OnceLock. See Server.
MLXCEL_REASONING_ALIAS_FIELD reasoning, none reasoning Native environment equivalent of --reasoning-alias-field. reasoning duplicates every Chat Completions reasoning_content value into an identical reasoning field; none suppresses only that duplicate alias. It does not change whether a model thinks, the existing reasoning_content field, Responses API reasoning events, or Anthropic thinking blocks.
MLXCEL_REASONING_ECHO_MAX_BYTES non-negative integer bytes; 0 disables 16777216 (16 MiB) Byte budget of the store that re-injects generated reasoning into content-only assistant history turns while the prompt cache is on (/v1/chat/completions, /v1/responses, /v1/messages), so a template re-renders an earlier turn the way it was generated and the follow-up still hits the prompt cache. The store also caps entries at 4096, refuses a single trace larger than an eighth of the budget, and evicts least-recently-used entries. An unparsable value warns and keeps the default. Read when the server starts.

Embedding and Reranking

Variable Values Default Notes
MLXCEL_EMBEDDING_MODEL local path or HuggingFace repo-id unset Embedding checkpoint loaded beside the primary model. LLAMA_ARG_EMBEDDING_MODEL is the compatibility alias; --embedding-model takes precedence.
MLXCEL_RERANKER_MODEL local path or HuggingFace repo-id unset Reranker loaded beside the primary model. LLAMA_ARG_RERANKER_MODEL is the compatibility alias; --reranker-model takes precedence.
MLXCEL_IMAGE_MODEL local path or owner/name repo-id unset Image-generation checkpoint served on /v1/images/generations and /v1/images/edits beside the chat model in -m, resolved the same way. With no -m (and no LLAMA_ARG_MODEL or --hf-repo), it selects image-only serving: the checkpoint becomes the served model, chat and tokenizer routes answer 501, and /v1/models lists it with the image capability. --image-model takes precedence. A load failure is a startup error.
MLXCEL_EMBEDDING_BATCH_SIZE positive integer 16 Text inputs per embedding forward pass. --embedding-batch-size takes precedence.
MLXCEL_EMBEDDING_MAX_LENGTH positive integer tokens checkpoint-derived, capped at 8192 Lowers the derived embedding input limit. --embedding-max-length takes precedence.
MLXCEL_EMBEDDING_QUEUE_DEPTH positive integer 8 Bound on each embedding and reranking worker command queue. --embedding-queue-depth takes precedence.
MLXCEL_EMBEDDING_REQUEST_TIMEOUT_SECS unsigned integer seconds 120 Reply timeout used by embedding and reranking requests. 0 falls back to the default; the CLI flag takes precedence.
MLXCEL_RERANK_BATCH_SIZE unsigned integer 0 Pairs per reranker forward pass. 0 selects the kind default: 8 for text, 2 for multimodal.
MLXCEL_EMBEDDING_POOLING cls, mean, max, lasttoken checkpoint/family Diagnostic pooling override used when --pooling is absent. The CLI flag outranks this variable and checkpoint configuration.

See Embedding and Reranking APIs for model-loading forms, queue errors, and normalization controls.

Build-Time Variables

These are read by the mlxcel-core build script. CUDA builds also use non-MLXCEL_* variables such as CUDA_HOME and MLX_CUDA_ARCHITECTURES; see Building from Source.

Variable Values Default Notes
MLXCEL_BUILD_METAL 1/0, on/off, true/false, yes/no on for macOS builds Overrides the CMake MLX_BUILD_METAL setting. Invalid values fail the build.
MLXCEL_BUILD_ACCELERATE 1/0, on/off, true/false, yes/no on for macOS builds Overrides the CMake MLX_BUILD_ACCELERATE setting. 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 such as x86-64-v3 for binaries that run on machines other than the build host; none omits -march entirely.
MLXCEL_GIT_COMMIT any string, for example a commit hash unset Read at compile time by the mlxcel crate (option_env!), not by a build script: the value present in the cargo build environment is embedded as server.build.git_commit in the WebUI bootstrap response (GET /ui-api/v1/bootstrap). Unset reports null. Nothing in the repository sets it; packagers may. Setting it at run time has no effect.

ROCm builds (--features rocm) read ROCM_PATH (default /opt/rocm) for the ROCm CMake packages, hipcc, and rocminfo, and MLX_ROCM_ARCHITECTURES for the gfx targets (default: detected with rocminfo); see Building from Source. The variables the ROCm backend reads at run time are in ROCm Backend Variables.

Native Windows builds (experimental) read different variables. Windows CUDA builds read CUDA_PATH (CUDA_HOME is honored ahead of it) and CUDNN_PATH, the cuDNN root, only when cuDNN is not unpacked into the CUDA toolkit directory. Windows ROCm builds read ROCM_PATH, then HIP_PATH when ROCM_PATH is unset or empty (there is no default; the build fails naming both), and detect the gfx targets with the SDK's amdgpu-arch.exe instead of rocminfo; MLX_ROCM_ARCHITECTURES still overrides detection. On Windows these paths are needed only at build time. See Experimental Support.

OpenXLA builds with --features xla-iree read IREE_DIST, or one source-built runtime with its matching compiler: IREE_CUDA_HOME with IREE_CUDA_COMPILE, IREE_MACOS_HOME with IREE_MACOS_COMPILE, or IREE_ROCM_HOME with IREE_ROCM_COMPILE. scripts/iree/setup-cuda.sh, scripts/iree/setup-macos.sh, and scripts/iree/setup-rocm.sh (or make iree-rocm) produce those trees; see OpenXLA Backend Source Status.

MLX_BUILD_CUDA is a CMake option selected by the Cargo cuda feature, not a runtime switch. MLXCEL_MLX_COMMIT and MLXCEL_XLA_IREE_DIST are compile-time values emitted by build scripts and baked into the relevant crates. MLX_PIN_CMAKE_FILE is used only by the MLX-pin parser's fixture tests.

ROCm Backend Variables

A --features rocm build reads these variables in its ROCm backend (src/lib/mlxcelverse/overlays/rocm/mlx/backend/rocm/). They are tuning, A/B, and bisecting switches, not stable public API, and none is recommended for deployment. Unless a row says otherwise, a variable is read once per process, the first time its code path runs, so set it before starting the process. The table covers the operator-facing subset; the backend reads further kernel-selection switches whose names and defaults follow the ROCm overlay source. Build options and the current ROCm support status are in Building from Source.

Memory and Device

Variable Values Default Notes
MLX_ROCM_MAX_INFLIGHT_MB non-negative integer MiB; other values warn and keep the default 1024 Memory that committed but unfinished command batches may have allocated. A batch is committed once it has allocated a quarter of the budget, and the host waits for the oldest batch while more is in flight, so a prefill's transients are not all live at once (Meta-Llama-3.1-8B-Instruct-4bit peak drops from 20.60 GB to 6.14 GB). 0 turns the bound off. The memory estimator adds this budget as its own Backend in-flight term (backend_inflight_bytes in mlxcel inspect --json).
MLX_ROCM_GPU_WATCHDOG_SECS non-negative integer seconds; other values warn and leave it off 0 (off) Longest a single host wait for GPU work may block before it fails with an error. A GPU fault is reported without it in about a second; the watchdog is for a kernel that never finishes, so it must exceed the longest prefill the server accepts. A watchdog failure still requires restarting the process.
MLX_ROCM_FFT_CACHE_SIZE positive integer up to 2147483647; other values warn and keep the default 128 hipFFT plans the plan cache keeps. A smaller value trades memory for plan rebuilds.
MLX_ROCM_FORCE_WARP_SIZE 32 or 64; other values are ignored with a stderr line the device's wavefront size Overrides the host-side launch width. Debug only: a wrong value produces garbage output.
MLX_HSACO_CACHE_DIR directory path system temp dir under mlx/<version>/hsaco Root of the on-disk cache of JIT-compiled kernels; one subdirectory per GPU architecture.
ROCM_HOME, ROCM_PATH, HIP_PATH directory path; tried in that order, HIP_PATH on Windows only; an empty value counts as unset /opt/rocm when it exists (Linux); on Windows the include in or beside the ROCm runtime directory, then the SDK include recorded at build time when it still exists ROCm installation the runtime JIT (hiprtc) takes its headers from. With none set and no fallback, the first JIT compile fails. A value set at run time wins over the fallbacks and its headers are compiled into GPU kernels, so it must name a directory only trusted accounts can write. The AMD HIP SDK installer sets HIP_PATH system-wide, which then takes precedence over a runtime bundle's include: clear it for mlxcel, or point it at the bundle root.
MLXCEL_ROCM_RUNTIME_DIR absolute directory path (Windows only); an empty value counts as unset unset: the executable's directory when it holds the HIP runtime (amdhip64_7.dll for ROCm 10), else the SDK bin recorded at build time Directory every ROCm DLL is loaded from, normally the bin folder of an extracted mlxcel ROCm runtime bundle. When set it is the only candidate: a relative path, or a directory without the HIP runtime, stops the binary at the first ROCm call with a message instead of falling back, and a DLL missing from the chosen directory is never taken from another installation. The chosen directory is printed once as [mlx-rocm] ROCm runtime: <dir> (<source>). Code in the directory runs in the mlxcel process, so it must be writable only by trusted accounts. See Windows ROCm Runtime Bundle.

Quantized Matmul and MoE

Variable Values Default Notes
MLX_ROCM_GATHER_QMV_EXPERT_BATCHED 0 off, 1 on; other values keep the default; read on every call on Expert-batched gather_qmm kernel for sorted MoE prefill (bf16 or f16 affine 4- or 8-bit at group 64, and bf16 mxfp4 at group 32). 0 sends every scheme to the per-row gather kernels, which is the way to A/B or bisect a MoE regression.
MLX_ROCM_WMMA_QMM unset, 0, or 1 unset (fused WMMA kernel where the device has native WMMA and is not a low-CU iGPU) How quantized_matmul runs a bf16 affine GEMM with more than one row. 0 never uses the WMMA kernel; 1 uses it wherever the shape fits and ignores MLX_ROCM_WMMA_QMM_MAX_M.
MLX_ROCM_WMMA_QMM_MAX_M integer from 1 to 2147483647; other values warn and keep the default 128 on RDNA 3.5 (gfx1150 to gfx1152), no ceiling elsewhere Row count from which such a GEMM goes to dequantize plus hipBLASLt instead of the WMMA kernel.
MLX_ROCM_QMM_DEQUANT_CACHE_SIZE integer from 0 to 2147483647 0 (off) Number of dequantized weights the dequantize-plus-GEMM route keeps for reuse. The shapes that reach that route rarely hit a cache, so it is off; a workload that does hit can opt in with a positive size.
MLX_ROCM_QMM_DEQUANT_CACHE_MAX_BYTES bytes 268435456 (256 MiB) Byte bound on that cache when it is enabled.
MLX_NO_HIPBLASLT presence hipBLASLt on Takes the rocBLAS path for every GEMM instead of hipBLASLt, a diagnostic for suspected hipBLASLt numerics.

Defaults That Differ on ROCm

These variables have full rows elsewhere on this page; a ROCm build changes only their default.

Variable Default in a ROCm build Elsewhere
MLXCEL_CACHE_LIMIT 2GB unset
MLXCEL_FUSED_ADD_RMSNORM on off on Metal and CUDA
MLXCEL_FUSED_ROPE_APPEND on off on Metal and CUDA
MLXCEL_FUSED_MOE_SGY 2 8 on Metal and CUDA
MLXCEL_FUSED_MOE_MAX_DFF 8192 (the CUDA value) 4096 on Metal
MLX_MAX_OPS_PER_BUFFER 2000 see MLX Runtime Variables

Hardware and Kernel Diagnostics

These variables are for profiling, rollback, or experiments. Do not rely on them as normal deployment settings.

Variable Values Default Purpose
MLXCEL_NO_PADDED_PREFILL presence disables auto Disables M5+/Neural-Accelerator prefill tile alignment in the CLI and in mlxcel-server. Takes precedence over MLXCEL_FORCE_PADDED_PREFILL.
MLXCEL_FORCE_PADDED_PREFILL presence enables off Forces 32-token prefill tile alignment on any hardware, CLI and server, so the padded prefill and its cache trim can be tested without an M5 Neural Accelerator. Pure overhead off M5; for validation only.
MLXCEL_FORCE_PADDED_PREFILL_MASK presence enables off Forces an explicit padded prefill mask path.
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.
MLXCEL_TRACE_ARCH presence enables off Prints the CUDA architecture picture once per process to stderr at device init: 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 reports the quantized-matmul path the dispatcher picks on its first call (qmm_sm90, qmm_sm80, qmm_naive, qmv, or fp_qmv), which MLX exposes no other way. Diagnostic only, with no synchronization and no extra graph work; on a ROCm build it prints the device's gfx target and the compiled gfx list instead. On a Metal or CPU-only build it reports that no compute capability is available. An architecture mismatch is reported without this variable: a binary whose compiled architectures do not cover the host GPU refuses to start with an error naming both, which MLXCEL_DEVICE=cpu bypasses.
MLXCEL_V4_DENSE presence forces dense pooled attention off (sparse past index_topk) DeepSeek-V4 diagnostic. 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 that the sparse path approximates. Below index_topk this is already what runs, so it only changes long-context behavior. 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_V4_FLAT_INDEX presence forces the flat scan off (hierarchy when the size gate opens) DeepSeek-V4 diagnostic. Forces the sparse layers' pooled-row selection onto the flat O(Np) scan, disabling the two-stage HiSA hierarchy while leaving sparse attention engaged. Both paths score the same pooled rows the same way, so running one prompt each way separates a hierarchy that retained the wrong blocks from a wrong sparse combine, 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 materializes the [B, H, L, Np] score tensor the hierarchy exists to avoid.
MLXCEL_ENABLE_FUSED_QKV_SPLIT_ROPE presence enables off Enables an experimental fused QKV projection/split/RoPE path.
MLXCEL_GEMMA4_ENABLE_FUSED_QKV presence enables off Enables a Gemma 4 fused-QKV projection experiment.
MLXCEL_GEMMA4_NO_KV_SHARED_TAIL_TRIM 1 disables trim on Keeps the prefill at full query width through the KV-shared tail instead of cutting the hidden state, per-layer input and masks down to the sampled row at the first KV-shared layer. The trimmed branch is numerically neutral against the full-width one, so this exists to run both arms of a teacher-forced trace on one binary. Only *_last_* prefill entries trim, and only when num_kv_shared_layers > 0.
MLXCEL_CUDA_DISABLE_GEMV 1 enables; unset or any other value leaves it off off CUDA only. Diagnostic switch that sends dense matmuls MLX would route to its own gemv kernel (M == 1 with a transposed B, or N == 1, no bias) to cuBLASLt instead. It changes B=1 decode of bf16/f16 models; quantized matmuls and gather_mm keep their own kernels. Read once per process, and meant for the gemv versus cuBLASLt measurement rather than tuning.
MLXCEL_DSA_DENSE presence enables off Forces DeepSeek-V3.2/GLM-MoE-DSA to bypass the DSA indexer and use dense full attention for A/B diagnosis.
MLXCEL_MINIMAX_M3_DENSE presence enables off Forces MiniMax-M3 to bypass its block-sparse indexer and use dense attention for A/B diagnosis.
MLXCEL_FUSED_QK_NORM 1/true/on/yes enables; any other value or unset disables off (opt-in) Fused single-token QKV projection + Q/K RMSNorm + RoPE decode kernel for Qwen3 and Qwen3-MoE. Matches the graph path within RMS < 5e-3, but is not measured to speed up decode on any backend tested so far (M1 Ultra, M5 Max, GB10/CUDA), so it stays opt-in.
MLXCEL_FUSED_XIELU 0/false/off/no disables; any other value or unset enables on Fused single-launch xIELU kernel (Metal, ROCm) for the Apertus MLP activation. Greedy temp-0 decode is byte-identical to the elementwise path on Apple Silicon and measured faster on M1 Ultra and M5 Max; the ROCm port is byte-identical to the elementwise path on gfx1151. On CUDA and CPU the FFI falls back to an equivalent elementwise graph. Apertus only. Set to 0 to force the elementwise path.
MLXCEL_DISABLE_COMPILED_SWITCH_QGEGLU presence disables compiled path when supported Rolls back the Gemma 4 compiled Switch-QGeGLU decode path.
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. Only softcap families reach this path at all, so Gemma 2 (attn_logit_softcapping) is what it moves. Greedy output is byte-identical either way.
MLXCEL_DISABLE_SOFTCAP_GQA_DECODE_GROUPED 1 disables, 0 enables unset Legacy rollback/override for grouped softcap-GQA decode.
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_PIPELINE_GRANULARITY off, layer, block:N off Inserts layer-boundary async-eval hints for pipeline experiments.
MLXCEL_AUTOTUNE 1/true/on/yes to tune; cache/read/readonly/read-only to consume cached tactics only; 0/false/off/no or unset to disable off Shape-bucketed kernel autotuner. Tuning mode measures candidate tactics per shape bucket and caches the winners; read-only mode applies a cache produced elsewhere without measuring.
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.
MLXCEL_FUSED_MOE 0/false/off/no disable; any other value or unset enables on Fused single-token decode-MoE kernel, on Metal, CUDA, and ROCm (HIP port, verified on wave32 gfx1151 only). Experts wider than MLXCEL_FUSED_MOE_MAX_DFF stay on gather_qmm. Disable it when comparing a new MoE port against a reference implementation, since the fused path is not byte-identical on every checkpoint.
MLXCEL_FUSED_ADD_NORM 0/false/off/no disable; any other value or unset enables on Fuses a parallel-residual block's (attn + mlp) + x with the LayerNorm that consumes it into one kernel (Metal, ROCm) during decode, byte-identical to that backend's unfused pair. Falls back to the unfused pair on CUDA and CPU, above a 6656-wide row, or on mixed dtypes. Used by Cohere2.
MLXCEL_DECODE_MB_PER_BUFFER 0/off/false/no disable; positive integer 1000 on pre-M5 Apple Silicon, off elsewhere Metal command-buffer input budget applied to pipelined decode steps only (the generate loops and the server's lookahead decode), leaving prefill and synchronous steps on the device default. In decode every token reads the whole weight set, so the default budget commits a buffer every one or two layers and idles the GPU at each boundary; measured decode gains on M1 Ultra range from flat to about 20% by model. Prefill stays on the default because the same value there roughly doubles peak memory on long prompts. Ignored when MLX_MAX_MB_PER_BUFFER is set.
MLXCEL_DEBUG_KERNEL_BACKEND presence enables off Diagnostic. Prints once which GPU backend mlxcel's custom kernels resolved to (metal, cuda, rocm, or none) and whether that backend has custom kernel ports or uses MLX graph fallbacks.
MLXCEL_SDPA_FALLBACK_MAX_QUERIES non-negative integer 32 CUDA only. A masked SDPA call with 2 to N query rows over a longer key sequence (the speculative verify shape) bypasses cuDNN and takes MLX's ops fallback, because cuDNN rebuilds its execution plan for every new key length. 0 restores upstream dispatch.
MLXCEL_SDPA_PLAN_BUCKET_MAX_QUERIES non-negative integer 0 (off) CUDA only, opt-in. Buckets the cuDNN plan-cache key for array-masked verify calls with 2 to N query rows, so one plan serves many rounds. Measured slower than the ops fallback at most widths on moderate contexts, so it is off; 32 enables it.
MLXCEL_SDPA_PLAN_BUCKET_MAX_MB non-negative integer (MiB) 64 CUDA only. Per-tensor ceiling on the zero-padded K/V copy the bucketing path makes when K and V are not a slice of one cache buffer. 0 disables the copy path.
MLXCEL_SDPA_PLAN_DEBUG 1 enables off CUDA only, diagnostic. Logs one line per cuDNN SDPA call with the fields that decide plan reuse and whether the call built or reused a plan. Per call per layer, so do not leave it on.
MLXCEL_FUSED_MOE_SGY 1-32 8 (Metal, CUDA) / 2 (ROCm) Simdgroups (Metal), warps (CUDA), or wavefronts (ROCm) per threadgroup for the fused decode-MoE kernel. Output does not depend on it. Tune per hardware.
MLXCEL_FUSED_MOE_MAX_DFF positive integer 4096 (Metal) / 8192 (CUDA, ROCm) Expert-intermediate (Dff) upper bound for the fused path. Above it the caller falls back to gather_qmm.
MLXCEL_FUSED_MOE_RELU2 presence enables off Squared-ReLU fused routed-expert path for nemotron-class experts.
MLXCEL_FUSED_MOE_PARITY_CHECK unset or 0 disables; nonzero integer enables off Diagnostic. In-situ per-call parity and determinism probe for the fused decode-MoE kernel. Use it to confirm a checkpoint's fused output tracks the reference.
MLXCEL_FUSED_MOE_PARITY_THRESHOLD float 0.05 Normalized-RMS deviation threshold for the MLXCEL_FUSED_MOE_PARITY_CHECK probe.
MLXCEL_INKLING_MOE_KERNEL 0 disables; any other value or unset enables on (Metal) Gates the fused Inkling MoE decode kernels on the backend with a port. The route kernel replaces the post-gate-matmul routing chain (sigmoid, correction bias, top-k, logsigmoid softmax, scaling, shared-expert gamma broadcast), and the down-combine kernel folds the routed affine q4/g64 down projection with the top-k weighted expert sum for batches of up to 8. A backend without a port takes the graph path, so 0 is only a rollback.
MLXCEL_FUSED_ADD_RMSNORM 0/false/off/no disable; 1/true/on/yes enable off (Metal, CUDA) / on (ROCm) Fused residual-add plus RMSNorm decode kernel, collapsing two passes over the hidden state into one.
MLXCEL_FUSED_ROPE_APPEND 0/false/off/no disable; 1/true/on/yes enable off (Metal, CUDA) / on (ROCm) Fused q/k RoPE plus KV-append-layout decode kernel.
MLXCEL_SPARSE_PAGED_ATTENTION 0/false/off/no disable; any other value or unset enables on Fused sparse decode via page indirection. When a request's live pages are a small fraction of its allocation, the kernel gathers only the selected rows.
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.
MLXCEL_SPARSE_PAGED_DUMP 1/true/on/yes enable off Diagnostic. Prints each request's selected pool rows for a fused sparse decode. Synchronizes with the device, so it is a debugging tool rather than a production switch.
MLXCEL_PAGED_DECODE_SPLITS positive integer unset (memory-budget ceiling) Pins the NumSplits token-split launch shape of the fused paged-attention decode kernel instead of deriving it from the memory budget.
MLXCEL_PAGED_ATTENTION_NATIVE 1/true/on/yes force native; 0/false/off/no force gather; unset defers to adaptive selection adaptive Forces one arm of the paged-attention decode dispatch for A/B measurement. The fused kernel runs on Metal, CUDA, and ROCm; ROCm follows the CUDA selection rule.
MLXCEL_PAGED_SLAB_BLOCKS positive integer blocks, used verbatim; 0 pins the pool default derived from --ctx-size and --parallel Paged pool slab size in blocks. The fused kernels read one contiguous buffer per side, so a layer whose rows outgrow its first slab returns to the gather path. The server logs the resolved value at startup as 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 as a side effect keeps the fused path unreachable 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.
MLXCEL_PAGED_V2_MIN_KV_TOKENS non-negative integer tokens 4096 Visible KV tokens a lone-request launch must reach before paged decode v2 is dispatched. The floor excludes the one measured loss, batch 1 at 1024 tokens. An unparseable value falls back to the default.
MLXCEL_PAGED_V2_MIN_KV_TOKENS_PER_REQUEST non-negative integer tokens 512 Visible KV tokens per request a multi-request launch must reach. Set below the smallest measured win so a prompt landing just under a nominal 1K is not declined without a reason.
MLXCEL_PAGED_ATTENTION_V2 1/true/on/yes enable; unset or anything else disables off Diagnostic. Makes the library-only v1 entry point try the fused v2 kernel first. It does not gate the server's production v2 decode, which the token floors above govern.
MLXCEL_PAGED_DECODE_V2_CHUNK positive integer pages unset Diagnostic. Pins pages_per_chunk for every paged decode v2 launch, overriding both the plan's occupancy heuristic and any autotuned value. The plan clamps the value into the feasible range.
MLXCEL_PAGED_DECODE_V2_TARGET_CTAS positive integer derived (Apple: gpu_core_count * 8, floored at 64; elsewhere 512) Diagnostic. Moves the occupancy target the v2 plan binary-searches its chunk size against, for hosts whose derived device parallelism is wrong.
MLXCEL_QMV_MULTIROW 0 disables; any other value or unset enables on CUDA only. Weight-amortizing multirow qmv path. With broadcast weights and a small row window, one weight load serves several rows.
MLXCEL_QMV_MULTIROW_MAX_ROWS 1-8 8 CUDA only. Row-window ceiling for the multirow qmv path, narrowing the dispatch window.
MLXCEL_QMM_TILE_M positive integer auto CUDA only. Pins the M-tile of the quantized matmul kernel instead of selecting it per shape.
MLXCEL_QMM_TILE_N, MLXCEL_QMM_TILE_K positive integer 128, max(64, group_size) CUDA only, experimental. Pins the other quantized-matmul CTA tile axes. Invalid layouts can make NVRTC compilation fail; use only for kernel tuning.
MLXCEL_QMM_FOLD_BATCH 0 disables; any other value or unset enables on CUDA only. Folds a batch that shares one 2D quantized weight into rows in QuantizedMatmul::eval_gpu, so batched decode reads the weight once instead of once per batch entry. Below MLXCEL_QMM_FOLD_MIN_ROWS the same qmv kernel runs over the folded rows; above it one GEMM reads the weight once. Read once at first use; 0 restores per-batch dispatch.
MLXCEL_QMM_FOLD_MIN_ROWS integer 2 to 8; anything else keeps the default 5 on compute capability 12.1, 8 elsewhere CUDA only. Rows from which a folded M == 1 decode batch of a large weight (see MLXCEL_QMM_FOLD_MIN_ELEMENTS) leaves qmv for the GEMM kernels. From this width a batched row is no longer bitwise equal to the same row decoded alone. Verify blocks and prefill pieces (M > 1) always keep 8. Read once at first use.
MLXCEL_QMM_FOLD_MIN_ELEMENTS positive integer 16777216 (2^24) CUDA only. Weight size N * K from which MLXCEL_QMM_FOLD_MIN_ROWS applies. A smaller weight keeps qmv through 8 rows on compute capability 12.1 and through 7 elsewhere. Read once at first use.
MLXCEL_QMM_NAIVE_TILE_N 64 or 128; any other value or unset leaves the selection to the shared-memory budget unset CUDA only, diagnostic. Pins the N tile of the qmm_naive kernel. Read on every dispatch, not cached, so one process can run both widths and compare them bitwise; each width compiles its own module.
MLXCEL_TRACE_QMM_TILE presence enables off CUDA only, diagnostic. Prints the selected qmm_naive tile with the compiled kernel's register count, spill, and achieved occupancy, once per distinct kernel, to stderr.
MLX_DEBUG_KERNELS 1 enables off Diagnostic. Requests MLX runtime-JIT kernel source dumps, including custom Turbo kernel launches.
MLXCEL_GATHER_QMM_GROUPED 0 disables; any other value or unset enables on CUDA only. Sorted MoE prefill fast path in GatherQMM::eval_gpu.
MLXCEL_GATHER_QMM_GROUPED_MIN_ROWS positive integer 8 CUDA only. Amortization threshold for MLXCEL_GATHER_QMM_GROUPED: the fast path activates once the sorted batch reaches this many rows.
MLXCEL_SSM_KERNEL 0 disables; any other value or unset enables on Fused single-token SSM decode kernel (ssm_update_kernel) for the hybrid SSM models (granite-4.0-h, falcon-h1, plamo-2, nemotron-h) on every backend with a port: Metal, CUDA, and ROCm. A backend without a port, or a CPU default device, takes the graph path. Set to 0 to force the graph path.
MLXCEL_SSM_CUDA_KERNEL 0 disables; any other value or unset enables on Alias of MLXCEL_SSM_KERNEL. Either variable set to 0 forces the graph path on every backend.
MLXCEL_MAMBA1_SCAN_KERNEL 0 disables; anything else or unset keeps the default on (Metal, ROCm; CUDA for Jamba) Fused kernel for the Mamba1 selective scan in Jamba, Mamba, and Falcon-Mamba, for prefill and decode. On Metal and ROCm the kernel keeps the state in float32, so results differ slightly from the graph scan; on CUDA, Jamba uses a port that rounds like the graph scan and is bit-identical to it. Measured on M1 Ultra (Jamba reasoning 3B 4-bit): prefill +45% to +54%, decode about +10%. Set to 0 to force the graph scan. Checked on every call.
MLXCEL_ONE_BIT_KERNEL 0/false/off/no disable; unset or any other value keeps the kernels on (Metal) Affine 1-bit checkpoints (bits: 1) run on mlxcel's own fused Metal qmv (fewer than 16 rows) and tiled qmm kernels. 0 forces the dequantize graph, which is the parity oracle and the path CUDA and ROCm always take. Read once per process.
MLXCEL_SDPA_DETERMINISTIC 1 (any non-zero integer) enables; words such as true read as off off (0) CUDA only. Makes MLX's one-row decode SDPA bitwise reproducible by routing it to the sdpa_vector kernel and keeping stream-K cuDNN plans out, at a decode cost of 0 to 13%. Required by mlxcel-engine-parity --expect-identical. It does not change the server's paged decode kernels.
MLXCEL_SDPA_VECTOR_LARGE_D 0/false/off/no disable; any other value or unset enables on CUDA only. Lets the CUDA supports_sdpa_vector check accept large head dimensions.
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 a single mx::compile graph.
MLXCEL_SHAPELESS_COMPILE_AUDIT any value except 0 enables off Audit mode, not a serving setting. Wraps every shapeless mlx::core::compile site in the C++ bridge with an eager oracle: each call also runs the uncompiled function and records per-site calls, input signatures, and mismatches for shapeless_compile_audit_report. Must be set before the first compiled function is built. Driven by scripts/audit_shapeless_compile.sh.
MLXCEL_SAMPLING_GUMBEL 0/false/off/no disable; any other value or unset enables on Softmax-free Gumbel-max categorical sampling kernel on Metal, CUDA, and ROCm. Draws from the logits directly, skipping the explicit softmax.
MLXCEL_SAMPLING_REJECTION 0/false/off/no disable; any other value or unset enables on Sorting-free top-k, top-p, and min-p sampling by dual-pivot rejection on Metal, CUDA, and ROCm, which avoids the full-vocabulary sort. Top-k together with top-p routes to it up to vocabulary 32768 (152064 on ROCm).
MLXCEL_NVFP4_DENSE_REPACK 1/true/on/yes forces the dense fallback off (direct transcode) Selects the dense fallback instead of the direct ModelOpt-triplet NVFP4 transcode. On ROCm, which has no native NVFP4 kernel, dense affine 4-bit conversion is already the default, so this changes nothing there.
MLXCEL_NVFP4_NATIVE_REPACK 1/true/on/yes selects native NVFP4 inside the dense fallback native by default Compatibility selector for the dense NVFP4 fallback path. Ignored on ROCm, which always converts ModelOpt NVFP4 to affine 4-bit; the load log says so.
MLXCEL_DISABLE_FUSED_GLOBAL_SCALE 1/true/on/yes disable off (fold on) Disables the NVFP4 global-scale fold, applying the scale as a separate step.
MLXCEL_CUDA_F16_NORMALIZE 1/true/on/yes enable; 0/false/off/no disable on below sm80, off (opt-in) at sm80 and later CUDA only. Load-time bf16 to f16 normalization of weights for the single-dtype decode graph. The default is split because the hardware is. Ampere and later have bf16 ALUs and gain nothing from f16, so they keep bf16. Volta and Turing (sm70, sm75) 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; there the conversion happens at load and quantized checkpoints are in scope, since their packed planes stay u32 and only the bf16 side-data moves. Set to 0 to keep bf16 on pre-Ampere.
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. Below sm_80 that list is down to apertus, whose xIELU squares its input and produces NaN at every scored position in f16. Use it only for a checkpoint whose activations you have reason to believe differ, or to take the measurement on a family the list does not yet know.
MLXCEL_CUDA_RUNTIME_DIR absolute directory path (Windows CUDA builds only); an empty value counts as unset unset: the executable's directory when it holds the CUDA runtime DLL, else the toolkit and cuDNN bin\<arch> recorded at build time Directory every CUDA, cuDNN, and NVRTC DLL is loaded from, normally the bin folder of an extracted mlxcel CUDA runtime bundle. When set it is the only candidate: a relative path, or a directory without the runtime DLL, stops the binary at its first CUDA call with a message instead of falling back. The chosen directory is printed once as [mlx-cuda] CUDA runtime: <dir> (<source>). Code in the directory runs in the mlxcel process, so keep it writable only by trusted accounts.
MLXCEL_CUDA_PINNED_LOAD 0/false/off disable; anything else or unset keeps the default on CUDA and ROCm (MLXCEL_ROCM_PINNED_LOAD on ROCm, with MLXCEL_CUDA_PINNED_LOAD also honored as a fallback). Uploads checkpoint tensors through four reused 64 MiB pinned host staging slots instead of one per-tensor buffer, so the read into one slot overlaps the copy out of the previous one. The 256 MiB of pinned memory is allocated on the first upload and kept for the process; if that allocation fails, one warning is printed and the pageable path is used. 0 forces the pageable path; the uploaded bytes are identical either way.
MLXCEL_ROCM_PINNED_LOAD 0/false/off disable; anything else or unset keeps the default on ROCm spelling of MLXCEL_CUDA_PINNED_LOAD, which is also honored as a fallback. Same pinned host staging behavior and the same pageable fallback.
MLXCEL_CUDA_GRAPH_CAPTURE_BARRIER 0/false/off/no disable; 1/true/on/yes, unset, or any other value enables on CUDA only. Makes every captured child graph inside an MLX CUDA graph (a cuBLASLt matmul, a cuDNN convolution or attention, any other stream capture) a barrier that joins every open branch before it and gates every node added after it, on top of the data dependencies MLX derives from buffers. Upstream orders a child graph only by its declared buffers, which on GB10 let unrelated kernels overlap and corrupted cuBLASLt f32 GEMMs in about 40% of fresh processes. Set to 0 for an A/B against upstream ordering.
MLXCEL_CUDA_GRAPHS_DOT_VERBOSE 1/true/on/yes enable; 0/false/off/no, unset, or any other value leave it off off CUDA only. With MLX's MLX_SAVE_CUDA_GRAPHS_DOT_FILE set, 1 dumps each graph with cudaGraphDebugDotFlagsVerbose, so kernel and memset node parameters (grid, shared memory, memset target) appear next to the node types; that is what shows which buffers concurrent nodes touch.
MLXCEL_QMV_ROWS_PER_WARP 1, 2 or 4; any other value 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 it read the whole activation vector, 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 activation traffic by R and leaving weight traffic alone; output is byte-identical to R = 1 at temperature 0. 1 routes to the stock kernel and is the rollback. Read once per process.
MLXCEL_QMV_MIN_BLOCKS 1, 3 or 4; any other value 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 registers: for a 256-thread block, 3 caps at 85 registers per thread and 4 caps at 64. The budget can be set too tight as well as too loose, so this rather than the row count carries the tuning. 1 is the unconstrained case and is the rollback. Read once per process.
MLXCEL_KEEP_BF16 any value enables off Measurement instrument, not a supported configuration. Skips load-time bf16 to 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. It is not evidence that bf16 beats f16: on families with no forced-bf16 subset the two land within noise of each other on the same hardware. What costs is a checkpoint holding both dtypes, because MLX promotes f16 with bf16 to f32 at every boundary between them. Read per load.
MLXCEL_BAILING_LINEAR_CHUNKED_PREFILL 0/false/off/no disable; any other value or unset enables on Chunked closed-form prefill for the bailing_moe_linear (Ling / Ring) gated linear attention, evaluated over 64-token chunks.
MLXCEL_PROFILE_NEMOTRON_MOE presence enables off Diagnostic. Forces evaluation at the router, routed-expert, combine and shared-expert boundaries inside fused_moe_forward, then emits one timing line per boundary.

Logging, Profiling, and Capture

Most of these switches force synchronization or extra graph work and can change throughput measurements.

Variable Values Default Purpose
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 the ordinary SDPA path runs instead. Off-switch only: it cannot turn the route on where the hardware test 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. Read once per process; inert on non-Metal builds.
MLXCEL_TRACE_DTYPE presence enables off Prints intermediate dtypes inside the fused Mamba2 forward (after the input projection and after the SSM update), evaluating the graph at each print.
MLXCEL_FORCE_SYNC presence enables off Keeps every decode step synchronous: disables the server BatchScheduler's lookahead decode pipeline and the raw-completion client's B=1 pipeline, which mlxcel generate (including the plain rounds of --prompt-lookup), the inference session, mlxcel-bench-decode, and the parity harness arms decode on. Greedy output is identical either way.
MLXCEL_PROFILE_TTFT presence enables off Emits one [TTFT] line per mlxcel generate --profile generation and per mlxcel-bench-decode measured pass, breaking the pre-first-token phase into setup, build (lazy graph construction, host only), sample, eval (the single blocking evaluation, which also materializes every weight, loads every JIT module, and instantiates every CUDA graph) and post, alongside the reported prefill and the unattributed residual. The sampler graph build runs inside the first evaluation, so sample reads 0 and eval carries both.
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_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_METAL_CAPTURE_PATH path unset Starts a Metal capture around the timed completion of plain mlxcel generate (after its one-token warmup) and around mlxcel-bench-decode's measured pass; requires MTL_CAPTURE_ENABLED=1.
MLXCEL_BENCH_PHASE_MARKS exactly 1 (trimmed) off mlxcel-bench-decode only. Prints host-clock [phase] <label> monotonic_ns=... boottime_ns=... marks to stderr so a kernel trace (for example rocprofv3) can be cut to the measured decode; scripts/rocm_decode_profile.py reads them.
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.
MLXCEL_LANG_BIAS_COUNTERS any non-empty value other than 0 enables off Collects the pre-bias suppression counters mlxcel_lang_bias_tokens_suppressed_total and mlxcel_lang_bias_byte_fragment_suppressions_total. Reading the pre-bias argmax costs a device synchronization on every decode step of a request with a token bias (logit_bias, ignore_eos, --lang-bias), so those requests leave the batched fused sampler and the lookahead pipeline while this is set. Read once per process.
MLXCEL_VOICECHAT_PROFILE_STAGES any non-empty value other than 0 enables off Adds sub-stage timers (for example perception.layers, tts.backbone, codec.decoder_istft) to the Nemotron VoiceChat mlxcel generate --stream --profile summary. The timers force evaluation at every boundary, so that run's frame totals are not the real-time factor.
MLXCEL_VOICECHAT_PROFILE_FRAMES file path unset Writes every VoiceChat frame's timings, cold frames included, as a JSON array for steady-state analysis of long inputs.
MLXCEL_APC_TRACE 1 enables (exact match only) off Prompt-cache trace logging, recording per-reason cache reject counts and hit/miss decisions. Unlike presence-based switches this matches 1 exactly.
MLXCEL_API_KEY secret string unset Not an API-key source: nothing in mlxcel authenticates with it (use --api-key or LLAMA_API_KEY). When set, its value is registered as a log secret and redacted from every log sink, like HF_TOKEN and LLAMA_API_KEY.

MLXCEL_PROFILE_PIPELINE, MLXCEL_PROFILE_PIPELINE_DETAIL, MLXCEL_EXPORT_DECODE_DOT, MLXCEL_CAPTURE_DECODE, MLXCEL_TRACE_ASTYPE, MLXCEL_TRACE_ENABLE_THINKING, MLXCEL_TRACE_BATCH_ROWS, MLXCEL_TRACE_IMAGE, MLXCEL_TRACE_IMAGE_PROMPT, MLXCEL_TRACE_LORA, MLXCEL_TRACE_LORA_MODE, MLXCEL_TRACE_TP_SIZE, and MLXCEL_LOGIT_TRACE_SCALAR_REFERENCE are not read by the current code. For a Metal trace of decode, use MLXCEL_METAL_CAPTURE_PATH; for pre-first-token timing, use MLXCEL_PROFILE_TTFT.

Test and CI Variables

These are intended for repository automation rather than normal end-user operation. See Testing.

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 test-serialization guard to a warning. The guard otherwise fails when the mlxcel-core suite is about to run with more than one test thread under the cuda feature.
MLXCEL_ALLOW_CONCURRENT_GPU_TESTS Downgrades the GPU-exclusivity guard to a warning. The guard otherwise fails when a second mlxcel-core test binary is already running, because two suites on one device interfere.
MLXCEL_TEST_VIDEO Turns the graceful skip in an ffmpeg-backed video test into a hard failure. Only the exact string 1 enables it. make verify-test-video sets it after installing ffmpeg.
MLXCEL_REQUIRE_PINNED_CHECKPOINTS Turns the graceful skip in a pinned-checkpoint contract test into a hard failure. Only the exact string 1 enables it. Meant for the machine that owns the pinned Muse Glimmer checkpoint, where an unconditional skip would silently disable that checkpoint's only coverage.
MLXCEL_REQUIRE_MODELS Turns a missing real checkpoint into a test failure instead of a skip, in tests/common/mod.rs and the real-model parity suites under src/; any value except 0 enables it. The real-checkpoint WebUI Makefile gates refuse to run without MLXCEL_REQUIRE_MODELS=1.
MLXCEL_AUDIO_PROBE_DIR Writes audio-encoder parity probes into the named directory.
MLXCEL_SPEC_DIST_SAMPLES Overrides the sample count of speculative-decoding distribution tests.
MLXCEL_TEST_GUMBEL_SAMPLES Overrides the Gumbel goodness-of-fit test sample count.
MLXCEL_TEST_REJECTION_SAMPLES Overrides the rejection-sampling test sample count.
MLXCEL_DUMP_CONFIG, MLXCEL_DUMP_OUT, MLXCEL_DUMP_DIR Inputs and outputs for ignored OpenXLA StableHLO dump tests.
MLXCEL_MOE_N Row count for the ignored OpenXLA MoE probe; defaults to 4.
MLXCEL_FREEZE_GOLDENS, MLXCEL_FREEZE_MOE Exact value 1 allows retained OpenXLA tests to rewrite checked-in golden assets. Never set these in a normal test run.

Not Environment Knobs

Some MLXCEL_-looking strings in the tree are intentionally omitted:

  • MLXCEL_BENCH in shell scripts is a local script variable, not a user-facing environment setting.
  • MLXCEL_KEY is an internal Rust constant name that stores the real key MLXCEL_PROMPT_CACHE_ENABLED.
  • __MLXCEL_MINICPMO_IMAGE__ is an internal prompt sentinel.
  • MLXCEL_CUDA_ARCHITECTURES and MLXCEL_ROCM_ARCHITECTURES are compile-time values build.rs hands to rustc, and MLXCEL_CUDA_INCLUDE_DIR, MLXCEL_ROCM_INCLUDE_DIR, and MLXCEL_ROCM_BIN_DIR are compile definitions set by the MLX CMake build. None is read from the shell at run time; use MLX_CUDA_ARCHITECTURES and MLX_ROCM_ARCHITECTURES to choose targets.
  • MLXCEL_TURBO_KERNEL_* names in CMake files are CMake list variables, not process environment variables.
  • MLXCEL_BRIDGE_METAL_BACKEND, MLXCEL_BRIDGE_CUDA_BACKEND, MLXCEL_BRIDGE_ROCM_BACKEND, and MLXCEL_BRIDGE_GPU_BACKEND are C++ preprocessor definitions selected by build.rs, not shell environment settings. Bridge C++ that calls MLX GPU helpers sits behind MLXCEL_BRIDGE_GPU_BACKEND so CPU-only builds still link.
  • LLAMA_ARG_LANG_BIAS_POLICY and LLAMA_ARG_LANG_BIAS_CONFIG appear only in a warning that explicitly states they are unsupported; use --lang-bias-policy and --lang-bias-config.
  • MLX_VLM_DRAFT_KIND appears in a legacy error hint but is not read; use --draft-kind or the documented LLAMA_ARG_DRAFT_KIND / MLXCEL_DRAFT_KIND aliases.
  • MLX_DISABLE_COMPILE is MLX's own switch that turns off graph compilation (compile). mlxcel's Rust code does not parse it; MLX, including the ROCm backend, reads it.
  • MLX_API, MLX_GET_TYPE, MLX_MIXED_BINARY_OP, MLX_MTL_PRAGMA_UNROLL, MLX_CUDA_SM90A_ENABLED, and MLX_CCCL_DIR are C/C++ compile-time macros, not shell settings. The similarly named runtime header override is MLXCEL_CCCL_DIR.
  • Historical investigation notes may mention old experiment names such as MLXCEL_DISABLE_GEN_STREAM or MLXCEL_ENABLE_COMPILED_SWITCH_QGEGLU; this revision does not read those names in executable code.