Two thread-safety races in vendored mlx 0.30.6 Metal backend under concurrent dispatch (encoder lifecycle + cross-stream signal-event)

[Ogilthorp3]

Hi! Filing this in the spirit of being helpful — we've shipped two PRs to mlx-rs this month (#347 + #348) and want to keep contributing useful signal. We hit two distinct Metal-backend race conditions in vendored mlx 0.30.6 while running concurrent inference on Apple Silicon, and both are reproducible in seconds under modest concurrency. Wanted to share what we learned.

Setup

• Mac mini M4 (128 GB unified)
• macOS 25.4.0 Darwin
• mlx-rs 0.30.6 (vendored via mlx-sys submodule)
• Rust binary: axum HTTP server with mTLS, OpenAI-compatible chat-completions endpoint
• Workload: Qwen3.6-35B-A3B-4bit MoE LLM + Qwen3-VL vision tower
• Concurrency model: two pipelines, each with its own tokio::sync::Mutex (Mutex<LoadedModel> for LM, Mutex<VisionForward> for vision). Each pipeline serializes within itself; the two pipelines do not synchronize across each other.

Bug 1 — encoder-lifecycle race when sharing the default stream

When vision and LM dispatch concurrently on the same default stream (per-pipeline mutex doesn't synchronize across pipelines), the encoder lifecycle races between mlx::core::metal::Device::get_command_encoder and CommandEncoder::~CommandEncoder. We've captured two crash signal variants depending on which thread wins the race:

SIGABRT (5 reproductions over 2026-05-08):

-[IOGPUMetalCommandBuffer setCurrentCommandEncoder.cold.1]:??
backtrace:
  mlx::core::fast::LayerNorm::eval_gpu
    → mlx::core::metal::Device::get_command_encoder
    → -[IOGPUMetalCommandBuffer setCurrentCommandEncoder]

SIGSEGV (2026-05-09 23:51:04):

crash site: mlx::core::binary_op_gpu
  → CommandEncoder::dispatch_threads
  → use-after-free in std::unordered_map<MTL::Resource*, void*>
  (resource map in mlx/backend/metal/device.cpp)

Both fire only when multimodal pipeline interleaves with another concurrent request. Both go through the same encoder lifecycle code in mlx/backend/metal/device.cpp. The CommandEncoder destructor does call endEncoding() + release(), but the cached encoder lookup in Device::get_command_encoder(int index) plus the stream.encoder = nullptr in Device::end_encoding(int index) doesn't appear safe under concurrent access from different threads.

Bug 2 — cross-stream signal-event-on-uncommitted-encoder race

Naturally we tried the architectural fix: give each pipeline its own mlx Stream so encoders are no longer shared. We wrapped each GPU-dispatch site in:

with_new_default_stream(MlxStream::new_with_device(&Device::gpu()), || {
    // dispatch ops here
})

This created independent DeviceStream per pipeline → independent encoder per pipeline → no shared resource. Bug 1 went away. We exposed Bug 2:

-[IOGPUMetalCommandBuffer encodeSignalEvent:value:]:427: failed assertion 
   `encodeSignalEvent:value: with uncommitted encoder'

Reproduces in <1 second under 8-way concurrent load (4 vision workers + 4 text workers). When one stream's output is consumed by another (e.g., vision tower output flowing into LM decode), mlx encodes a Metal signal-event so stream B can wait on stream A. The signal-encoding requires a committed encoder, and the cross-stream sync path doesn't always commit before signaling under concurrent dispatch.

Reproducer

5-minute concurrent stress harness, MIT-licensed, ~200 lines of bash:

https://github.com/Ogilthorp3/Claude_Code/blob/main/tools/cathedral/concurrent-stress.sh

Defaults: 5-min duration, 4 workers per pipeline (8 total), curl with mTLS to 127.0.0.1:1337. Watches /Library/Logs/DiagnosticReports/sanctum-mlx-*.ips for new crash reports during the run. Returns pass/fail.

Pre-fix runs (Bug 1 chasing): hourly SIGABRT/SIGSEGV under canary + guardian probe interleaving (~60s + 60s cadences crossing).
Distinct-streams attempt (Bug 2): FAIL within 1 second of stress.
Global-mutex workaround (see below): 353 requests, 0 errors, 0 crashes over 5 min.

Workaround we're shipping

Application-layer global Mutex<()> in our AppState, acquired at the top of every chat-completion handler. Net: one inference task in flight globally, regardless of pipeline. Avoids both bugs because only one stream is ever active.

Trade: serialized inference (~850 ms p50 under our workload). Fine for our 1-2 concurrent peak; bad for high-throughput. Streaming handlers deferred (the spawn-then-return pattern needs lock_owned on Arc<Mutex>, and our canary + guardian don't exercise streaming).

Merged: https://github.com/Ogilthorp3/sanctum-rs/pull/10

What would help us most

In rough priority order — happy to contribute on any of these:

1. Documentation note on mlx_rs::Stream and Device thread-safety guarantees. Even a one-liner ("Stream operations on the shared default stream require external synchronization") would have saved us from finding both bugs the hard way.
2. Safer Rust wrappers that hold the right granularity of lock internally — perhaps a SyncStream newtype that wraps Stream + an internal mutex. Optional opt-in, no perf cost for users who don't need it.
3. Underlying mlx C++ fix if upstream (ml-explore/mlx) is willing — happy to relay the report there with the IPS files attached.
4. Stream API ergonomics: Stream::new() actually returns the default stream (via mlx_get_default_stream), not a new one. The actual distinct-stream constructor is Stream::new_with_device(&device) → mlx_stream_new_device. The naming was a footgun for us; a doc note or rename would prevent the same confusion downstream.

Happy to provide additional repros, IPS files, or contribute fixes if you can point us at the right level. Thanks for mlx-rs — it's the foundation of everything we're building on Apple Silicon.