SIGSEGV in QuantizedMatmul::eval_gpu during long token generation on Mac Studio M2 Ultra

[Thump604]

Description

Sporadic SIGSEGV crashes during quantized inference on Mac Studio M2 Ultra (Mac14,14). The crash occurs in Apple's Metal driver (AGXMetalG14X) when MLX's QuantizedMatmul::eval_gpu creates a new command encoder while the driver is coalescing a previous compute command buffer.

The crash is not caused by concurrent API requests — we serialize all inference requests through an asyncio lock (one request at a time). It appears to be an internal race condition between MLX's GPU dispatch threads and Metal's command buffer management.

Environment

• Hardware: Mac Studio M2 Ultra (76-core GPU, 128GB unified memory, wired to 115GB via MLX_MAX_TBP_SIZE_GB)
• macOS: 26.3 (25D125)
• MLX: 0.31.0
• Python: 3.12.12
• Model: Qwen3.5-27B-8bit (MLX quantized format)
• Serving: vllm-mlx 0.2.6

Reproduction

Crashes happen during long token generation (thousands of tokens, e.g. generating large JSON output). Short responses (< 100 tokens) work reliably. 37 crashes collected over 2 days (March 13-14) during normal single-request use. The server auto-restarts via LaunchAgent after each crash.

# Typical invocation (via vllm-mlx serve)
vllm-mlx serve ~/ai-models/mlx_models/Qwen3.5-27B-8bit \
    --host 127.0.0.1 --port 8080 \
    --tool-call-parser qwen3_coder \
    --reasoning-parser qwen3 \
    --enable-auto-tool-choice --mllm --timeout 2400

Crash Details

Exception Type:  EXC_BAD_ACCESS (SIGSEGV)
Exception Subtype: KERN_INVALID_ADDRESS at 0x0065c62f691dadd0 -> 0x0000462f691dadd0
                   (possible pointer authentication failure)

Faulting Thread: 21 (of 52 total)
Threads with MLX frames: [1, 2, 3, 4, 5, 6, 7, 8, 9, 21, 51] (11 total)

Faulting Thread Stack

0: libobjc.A.dylib :: objc_msgSend
1: libobjc.A.dylib :: -[NSObject isMemberOfClass:]
2: AGXMetalG14X :: -[AGXG14XFamilyCommandBuffer tryCoalescingPreviousComputeCommandEncoderWithConfig:nextEncoderClass:]
3: AGXMetalG14X :: -[AGXG14XFamilyCommandBuffer computeCommandEncoderWithConfig:]
4: AGXMetalG14X :: -[AGXG14XFamilyCommandBuffer computeCommandEncoderWithDispatchType:]
5: libmlx.dylib :: mlx::core::metal::CommandEncoder::CommandEncoder(mlx::core::metal::DeviceStream&)
6: libmlx.dylib :: mlx::core::metal::Device::get_command_encoder(int)
7: libmlx.dylib :: mlx::core::qmv(...)
8: libmlx.dylib :: mlx::core::QuantizedMatmul::eval_gpu(...)
9: libmlx.dylib :: mlx::core::gpu::eval(mlx::core::array&)
10: libmlx.dylib :: mlx::core::eval_impl(...)
11: libmlx.dylib :: mlx::core::async_eval(...)

A second crash on the same day showed KERN_INVALID_ADDRESS at 0x0000000000000000 (null pointer dereference) on a different thread (28), also originating from MLX GPU evaluation.

Mitigations Attempted (all failed)

Mitigation	Result
MLX_NUM_THREADS=1	Still crashes
TOKENIZERS_PARALLELISM=false	Still crashes (eliminated 24 idle rayon threads)
External request serialization (asyncio Lock)	Already in place — crash is internal

Thread Analysis

At crash time, 52 threads were active in the process:

Threads	Role	State
T1-T9	MLX ThreadPool workers	Idle (condition_variable::wait)
T10-T20	Python eval frames	Running
T21	MLX GPU dispatch (CRASHED)	QuantizedMatmul::eval_gpu → Metal tryCoalescingPreviousComputeCommandEncoder
T22-T45	HuggingFace tokenizers rayon	Idle (sleeping)
T46	IOGPU	_submitCommandBuffers
T51	MLX eval_impl	async_eval → waiting

The key observation: T21 and T46 were both active in GPU command buffer operations simultaneously — T21 creating a new command encoder while T46 submitted command buffers. The pointer authentication failure suggests a use-after-free or stale reference in the Metal command buffer being coalesced.

Additional Context

• We also see this in mlx-examples #866 where M2 Ultra fails while M3 Max succeeds on the same model.
• The crash only occurs during extended quantized matmul operations (long token generation). Short generations are stable.
• Full macOS crash reports (.ips files) are available if needed.

[datawakes]

Reproducing on M3 Ultra / macOS 26.3 + 26.3.1 / oMLX

Same crash here. Adding another data point:

• Hardware: Mac Studio M3 Ultra (80-core GPU, 256GB unified memory)
• macOS: 26.3 (25D125) and 26.3.1 (25D2128) — both affected
• MLX: 0.31.0 (bundled in oMLX 0.2.7)
• Models: GLM-4.7-Flash-MLX-8bit, Qwen3.5-122B-A10B-MLX-9bit
• Serving: oMLX 0.2.7 (not vllm-mlx)

Crash backtrace looks identical -> objc_msgSend -> tryCoalescingPreviousComputeCommandEncoderWithConfig -> CommandEncoder::CommandEncoder -> Gather::eval_gpu. Getting EXC_BAD_ACCESS (KERN_INVALID_ADDRESS) with pointer authentication failure.

The crash is frequent; oMLX's server process dies every few minutes under normal load (single sequential requests, not concurrent). Clearing the SSD KV cache did not help. No crashes observed on macOS 26.2.x with the same hardware and MLX version.

[zcbenz]

Can you make a simple demo showing how you are using MLX in a multi-thread environment? Making multi-thread inference work is something that I think we can and should implement, and I have put some details in #3078 (comment).

[Thump604]

We're not using MLX from multiple threads. We can't, because it isn't safe to. The serving stack (vllm-mlx) serializes all inference through:

1. asyncio.Lock — only one generation runs at a time
2. ThreadPoolExecutor(max_workers=1) — prefill offloaded to a single background thread
3. Single mx.stream(generation_stream) — all generation steps run on one dedicated Metal stream

There are no concurrent mx.eval() or mx.async_eval() calls from user code. The threading visible in the crash report is internal to MLX and Metal:

• T1-T9: MLX's own ThreadPool workers (idle at crash time)
• T21: MLX eval_gpu → CommandEncoder constructor (crashed)
• T46: Metal's IOGPU → _submitCommandBuffers
• T51: MLX async_eval → waiting

The race appears to be between MLX's command encoder creation on T21 and Metal's command buffer submission on T46, both operating on the same MTLCommandBuffer. The pointer authentication failure in tryCoalescingPreviousComputeCommandEncoder suggests a use-after-free in the buffer being coalesced.

That said, proper concurrent inference support in MLX would be huge for us. Right now we're leaving a lot of the M2 Ultra's capacity on the table. With 128GB of unified memory and a 76-core GPU, we have the headroom to serve multiple requests in parallel, but the single-threaded serialization means we're bound to one-at-a-time throughput. Even level 2-ii from your comment on #3078, independent arrays on separate streams that can be safely eval'd on any thread, would unlock parallel prefill + decode, which is the main bottleneck for serving workloads on Apple Silicon.

We can put together a standalone repro script for this crash if that would help, but the key point is this isn't a user-level threading issue. It's happening inside MLX's own GPU dispatch even with fully serialized inference.

[zcbenz]

The most confusing part to me is that async_eval and eval_gpu ran in 2 threads, because eval_gpu() is only invoked by async_eval and eval in the same thread, i.e. if there is a eval_gpu running in T21 then you must have invoked async_eval or eval in that thread.

Can you check if all async_eval/eval calls actually run in the same thread?

[Thump604]

Instrumented the running server by monkey-patching mx.eval and mx.async_eval to log threading.get_ident() and thread name on every call. Ran two requests (short prompt + 4096 token generation, long 4500-token prompt + 2048 token generation) and captured the report on shutdown:

=== EVAL THREAD REPORT (mlx#3216) ===
Total calls: 3077 | Unique threads: 2
  Thread 67162632192 (asyncio_0): async_eval x3074
  Thread 67162632192 (asyncio_0): eval x2
  Thread 8738812160 (MainThread): eval x1
Thread transitions: 1

All inference eval calls run on a single thread (asyncio_0). The one MainThread eval call happens during model loading at startup, before the event loop starts. Zero thread transitions during generation.

(The serving path uses SimpleEngine, which runs everything on the asyncio event loop thread. The ThreadPoolExecutor pattern mentioned earlier is in EngineCore, a different engine mode not in use here.)

This confirms eval_gpu is only invoked from one thread at a time. Yet the crash report shows T21 and T51 both inside MLX eval code simultaneously. @datawakes reproduced the same crash on M3 Ultra with a completely different serving stack (oMLX), and only on macOS 26.3, not 26.2.x. That pattern points to something internal to MLX's interaction with the Metal driver on 26.3.

Full .ips crash reports are available on request.

[Thump604]

Additional data point on crash frequency: 37 crash reports collected over March 13-14, all during normal single-request inference on Qwen3.5-27B-8bit. All crashes are in MLX's Metal interaction:

Signature	Count	Exception
tryCoalescingPreviousComputeCommandEncoder / CommandEncoder / eval_gpu	15	EXC_BAD_ACCESS (SIGSEGV) + EXC_CRASH (SIGABRT)
Other MLX/Metal frames (libmlx + AGXMetalG14X)	22	Same mix

All 37 on macOS 26.3 / 26.3.1, M2 Ultra. The process restarts automatically via LaunchAgent after each crash, so the high count reflects continuous use, not a single session.

Crash reports (.ips files, ~66-76KB each) available if useful for diagnosing the Metal driver interaction.

[Thump604]

Same crash just now with Nemotron-3-Super-120B-A12B (hybrid Mamba-2 + Attention + LatentMoE architecture), not just Qwen:

Exception: EXC_CRASH / SIGABRT
Faulting thread: 16

Metal :: MTLReportFailure
AGXMetalG14X :: -[AGXG14XFamilyCommandBuffer tryCoalescingPreviousComputeCommandEncoderWithConfig:nextEncoderClass:]
AGXMetalG14X :: -[AGXG14XFamilyCommandBuffer computeCommandEncoderWithConfig:]
AGXMetalG14X :: -[AGXG14XFamilyCommandBuffer computeCommandEncoderWithDispatchType:]
libmlx.dylib :: mlx::core::metal::CommandEncoder::CommandEncoder(mlx::core::metal::DeviceStream&)
libmlx.dylib :: mlx::core::metal::Device::get_command_encoder(int)
libmlx.dylib :: mlx::core::Arange::eval_gpu(...)

This time through Arange::eval_gpu, not QuantizedMatmul::eval_gpu. Same Metal driver path, different MLX primitive. macOS 26.3.1, MLX 0.31.1, M2 Ultra.

This plus @datawakes seeing it on M3 Ultra with oMLX confirms it's model-independent and serving-stack-independent. The common factor is macOS 26.3+ and the tryCoalescingPreviousComputeCommandEncoder code path in AGXMetalG14X.

[zcbenz]

Thanks for running the tests, it is probably a Metal bug, though there are still many confusing parts, for example CommandEncoder::CommandEncoder is only created when you create a stream for the first time, which should not happen in a long time running.

Can you check whether the crash ever happened in macOS 26.3.0?

Also, is there any assertion message printed in the terminal?

[Thump604]

macOS version: All 44 crash reports on this machine are macOS 26.3.1 (25D2128) — none on 26.3.0.

Here's the update timeline from softwareupdate --history:

Version	Installed	Crashes
26.2	Jan 13	0
26.3	Feb 18	0 (ran inference daily through Mar 11)
26.3.1	Mar 11	44 (first crash Mar 12, one day after update)

So the regression window narrows to 26.3.1 specifically, not 26.3.0. We ran quantized inference continuously on 26.3 for three weeks with zero crashes.

Assertion messages: Yes — two distinct assertions print to stderr before the process dies:

-[AGXG14XFamilyCommandBuffer tryCoalescingPreviousComputeCommandEncoderWithConfig:nextEncoderClass:]:1094: failed assertion `A command encoder is already encoding to this command buffer'

(5 occurrences)

-[_MTLCommandBuffer addCompletedHandler:]:1011: failed assertion `Completed handler provided after commit call'

(1 occurrence)

The first assertion suggests the issue isn't CommandEncoder::CommandEncoder at stream creation — it's the coalescing path (tryCoalescingPreviousComputeCommandEncoder) firing while an encoder is already active on the same command buffer. The second suggests a command buffer lifecycle issue where a completion handler is registered after the buffer has already been committed.

Crash breakdown: 28 EXC_CRASH (from the assertion aborts) + 13 EXC_BAD_ACCESS (PAC failures in the same Metal code path) across 44 reports. All on M2 Ultra, MLX 0.31.1, during long token generation (100–600+ seconds).