DeepSeek-V4 (Flash/Pro) on Apple Silicon: unbounded Metal residency growth during decode → [metal::malloc] Resource limit (499000) exceeded after ~11K tokens

Summary

On Apple Silicon, the DeepSeek‑V4 port (#1192) aborts generation with
RuntimeError: [metal::malloc] Resource limit (499000) exceeded after ~11,300 generated tokens, independent of prompt length. After the first failure the Metal command queue is wedged until a full process restart. This is the "looping / long‑context" failure several testers have hit in the #1192 thread. Root cause is an unbounded per‑decode‑step leak of live Metal buffers in the attention caches. A fix is in Blaizzy#25.

Environment

mlx‑lm @ Add DeepSeek-v4 (Flash/Pro) #1192 head (Blaizzy:pc/add-deepseekv4flash-model), mlx 0.31.2
Apple M4 Max, 128 GB, macOS 26.3 (also reproduced on M3 by others in the thread, ~4000 tokens)
mlx-community/DeepSeek-V4-Flash-2bit-DQ

Reproduce

Server (mlx_lm.server --max-tokens 65536), then stream a long generation — it breaks at ~11,314 tokens:

import json, urllib.request
body = {"model": MODEL, "stream": True, "max_tokens": 20000, "temperature": 0.0,
        "messages": [{"role":"user","content":"Count from 1 to 20000, one number word per line."}]}
# ... stream and count delta tokens; the connection drops at ~11,314 with the server-side OOM

In‑process (no server, converges in <1 min) — measures the leak rate directly:

import mlx.core as mx
from mlx_lm import load
from mlx_lm.models.cache import make_prompt_cache
model, tok = load(MODEL); cache = make_prompt_cache(model); MB = 1024**2
y = mx.argmax(model(mx.array(tok.encode("Count upward:\n"))[None], cache=cache)[:, -1], -1); mx.eval(y)
base = mx.get_active_memory() / MB
for s in range(500):
    y = mx.argmax(model(y[:, None], cache=cache)[:, -1], -1); mx.eval(y)
    if s % 100 == 0: print(s, f"{mx.get_active_memory()/MB - base:+.1f} MB")
# active memory grows ~linearly, no plateau (~0.2 MB/step). Adding
# `mx.eval([c.state for c in cache])` after the per-step eval flattens it to ~0.

Root cause

resource_limit (499000) is a count of live resident Metal buffers, not bytes (num_resources_ vs resource_limit_ in mlx/backend/metal/allocator.h, backed by the ResidencySet). No byte‑budget knob (set_memory_limit/set_wired_limit/set_cache_limit) affects it. At the OOM point only ~2–3 GB is leaked — it's the count that's exhausted.
The attention caches build un‑detached per‑step graphs. PoolingCache.update_and_fetch does self.pooled = mx.concatenate([self.pooled, px], axis=1) each step; pooled_N references pooled_{N-1} in its input graph, and the cache permanently holds the head, so every prior step's intermediate array — and its Metal buffer — stays resident. RotatingKVCache._update_in_place and the Batch* variants chain identically via sliced functional assignment. That's ~one retained buffer per layer per step → 499000 / 43 layers ≈ 11.3K steps, matching the observed failure point.
Realizing the network output (mx.eval(token)) does not detach those intermediate cache arrays; explicitly evaluating the cache state each step collapses the growth (~205 → ~7 KB/step). Standard in‑place KVCache writes don't hit this, which is why other models don't.

It is not a memory‑size problem, not a single‑forward‑pass / per‑command‑buffer problem, and not prompt‑cache‑length driven (a 58‑token prompt OOMs at the same step count).

Full analysis + the diagnostic path (ablation, force‑eval, ruling out mx.compile/RoPE/expert‑paging): full writeup. PR with the fix: Blaizzy#25.

docs(deepseek-v4): record upstream submission + add 300-req knowledge…

fix: materialize DeepSeek-V4 cache arrays

fix(M3): eval m3_idx_k after concat to break lazy-graph buffer leak

`ArraysCache.advance()` leaks one Metal buffer handle per token per un-drained cache

mlx_lm.server dies with [metal::malloc] Resource limit (499000) exceeded during decode on
hybrid/linear-attention models. resource_limit is a COUNT of live Metal resources, not bytes —
memory is flat while it happens.

Environment. M5 Max, 128 GB, macOS 27.0, Python 3.14, mlx 0.31.2 (0.32.1 identical),
mlx-lm 0.31.3, mlx-community/Qwen3.8-27B-{4,8}bit (qwen3_5, 64 layers, 48 linear).

The mechanism

mlx_lm/models/cache.py:685-689:

def advance(self, N):
    if self.lengths is not None:
        self.lengths -= N
    if self.left_padding is not None:
        self.left_padding -= N

N is a Python int, so -= N materialises a fresh scalar mx.array operand eagerly, then
builds an unevaluated subtract node on top of it. The node keeps that operand reachable, so it is
never freed, and mx.clear_cache() cannot reclaim it (it drains the free pool, and these are
live). One retained 4-byte buffer per call.

Isolated without any model (mlx 0.31.2, 20,000 iterations):

pattern	bytes/step
`x = x - 1`, never evaluated	4.0
`x = x - one` (pre-evaluated array operand), never evaluated	0.0
`x = x - 1`, `mx.eval(x)` each step	0.0

Graph depth alone costs nothing; the retained eager operand is the leak. Both conditions are
necessary. An unevaluated x = x - 1 chain raises the byte-for-byte identical error at exactly
499,000 iterations.

Why only the batch/server path

left_padding is None unless something sets it, and both sites that set it are batch-only:

cache.py:604 — ArraysCache.__init__(size, left_padding=None) sets it when the caller
passes one; make_cache in the plain path never does
generate.py:848 — c.left_padding = mx.array(left_padding) in _make_cache, the batch
cache builder, applied to the ArraysCache instances the model returns
cache.py:709 — cache.left_padding = mx.array([0] * B) in ArraysCache.merge, on the
all-empty early-return branch, reached via GenerationBatch.__init__ → generate.py:1036 _merge_caches(...)

(The other left_padding = mx.array(...) sites at cache.py:936, 973, 1140, 1278 belong to
BatchKVCache / BatchRotatingKVCache, which are not affected — see the asymmetry below.)

In stream_generate, make_cache builds ArraysCache(size=2) with no left_padding, so
advance() is a total no-op. Measured: plain decode is flat over 16,000 tokens; batch decode
leaks from token one.

For a single-sequence request max_padding == 0, so finalize() — which would set
left_padding = None (cache.py:681-683) — is skipped by the guard at generate.py:1150/1166,
and the array stays live for the whole completion.

Why only hybrid architectures

advance() is defined on ArraysCache and on no other cache class, and is invoked only from
linear/SSM layers (qwen3_5.py:198, inside GatedDeltaNet.__call__). A pure-attention model of
any depth constructs zero ArraysCache and executes zero advance() calls. Measured: a
48-layer pure-attention control at matched depth ran flat past 20,000 tokens.

Why the rate is 47, not 48

All 48 linear layers call advance(), but only one left_padding is ever read:
qwen3_5.py:251 sets ssm_idx = 0, :269 builds the mask from cache[self.ssm_idx] alone, and
:271-273 reuses that single mask for every linear layer. create_ssm_mask → make_mask
(cache.py:691-694) is the only reader. Cache 0's chain is therefore drained one step in arrears
through the logits; the other 47 accumulate.

47 handles/token against 499,000 less a ~2,100 baseline predicts a 10,572-token ceiling.
Measured: 10,560, three runs (10,562 / 10,561 / 10,560). Agreement 0.12%.

Corroborating at a different layer count: Falcon-H1 with 24 ArraysCache crashed at 21,641
tokens; 499,000 / 23 = 21,696.

Why no existing eval catches it

ArraysCache.state (cache.py:624-626) returns self.cache only — left_padding and lengths
are outside it. So all four mx.eval([c.state for c in ...]) sites (generate.py:442, 584, 1161, 1169) miss them, and all four are prefill anyway. GenerationBatch._step
(generate.py:1320-1379) evaluates only mx.async_eval(self._next_tokens, self._next_logprobs, token_context) and mx.eval(inputs, self._current_logprobs). Nothing in the decode loop touches
cache metadata.

Note for anyone attempting the obvious fix: adding the eval of c.state to the decode loop does
not work, precisely because left_padding is not in state. We measured that arm — it leaks
identically to the control.

The clearest view of it — `to_batch_cache`, `generate.py:846-852`

def to_batch_cache(c):
    if type(c) is KVCache:
        return BatchKVCache(left_padding)        # purpose-built batch class
    elif isinstance(c, ArraysCache):
        c.left_padding = mx.array(left_padding)  # raw field assignment, returned as-is
        return c

A KVCache is CONVERTED to BatchKVCache, which includes left_padding in state and whose
rotating sibling pins it with mx.depends. An ArraysCache is not converted at all — it is
mutated in place and handed back. There is no BatchArraysCache. The class that acquires
per-step mutable batch metadata is the one class that never got a batch-aware implementation.

The asymmetry, stated as one sentence

ArraysCache is the only batch-capable cache class whose per-step mutable metadata is neither
included in state nor pinned with mx.depends. The sibling classes already have the guard:

BatchKVCache.state (cache.py:990-995) returns k, v, self.offset, self.left_padding
BatchRotatingKVCache (cache.py:1203-1204 and 1256-1257) does
self.keys = mx.depends(self.keys, (self.left_padding, self.offset)) under the literal comment
# Make sure left_padding and offset are evaluated

Suggested fix

Any of, in order of locality — all inside cache.py:

Include left_padding and lengths in ArraysCache.state, and add a decode-loop flush,
since today the decode loop evaluates no cache state at all.
Mirror the mx.depends guard the batch KV caches already use, pinning the metadata into a
graph that is forced with the token.
mx.eval(self.left_padding) inside advance() — simplest, at the cost of one eval per layer
per step rather than one batched eval.

We deployed the equivalent of (3) as a batched per-step eval of left_padding/lengths in
GenerationBatch._step. Measured over 6,000 tokens: headroom flat with no drop (control fell
450,000 → 150,000), and 22% faster wall-clock — the growing unevaluated graph costs time as
well as handles, which also explains the superlinear decode slowdown reported in these threads.
End-to-end, a single request now completes 14,000 tokens where 10,560 was previously fatal.

Scope

.advance( has call sites across models/ for: falcon_h1, granitemoehybrid, kimi_linear, lfm2,
lfm2_moe, mamba2, nemotron_h, plamo2, qwen3_next, qwen3_5. Every hybrid/SSM family served through
the batch path is exposed, scaled by its own linear-layer count. Fixing it in ArraysCache fixes
all of them at once.

I'm merging this to #1662.

Materialize cache arrays after each forward pass (Metal residency fix)

Navigation Menu

Repository navigation

DeepSeek-V4 (Flash/Pro) on Apple Silicon: unbounded Metal residency growth during decode → [metal::malloc] Resource limit (499000) exceeded after ~11K tokens #1332

Summary

Environment

Reproduce

Root cause

`ArraysCache.advance()` leaks one Metal buffer handle per token per un-drained cache

The mechanism

Why only the batch/server path

Why only hybrid architectures

Why the rate is 47, not 48

Why no existing eval catches it

The clearest view of it — `to_batch_cache`, `generate.py:846-852`

The asymmetry, stated as one sentence

Suggested fix

Scope

Metadata

Assignees

Labels

Type

Projects

Milestone

Relationships

Development

DeepSeek-V4 (Flash/Pro) on Apple Silicon: unbounded Metal residency growth during decode → [metal::malloc] Resource limit (499000) exceeded after ~11K tokens #1332

Description

Summary

Environment

Reproduce

Root cause

Activity

fp2s92d5pd-bit commented on Aug 25, 2026

ArraysCache.advance() leaks one Metal buffer handle per token per un-drained cache

The mechanism

Why only the batch/server path

Why only hybrid architectures

Why the rate is 47, not 48

Why no existing eval catches it

The clearest view of it — to_batch_cache, generate.py:846-852

The asymmetry, stated as one sentence

Suggested fix

Scope

zcbenz commented on Aug 27, 2026

Metadata

Metadata

Assignees

Labels

Type

Projects

Milestone

Relationships

Development

Issue actions

`ArraysCache.advance()` leaks one Metal buffer handle per token per un-drained cache

The clearest view of it — `to_batch_cache`, `generate.py:846-852`