Skip to content

Fix a deadlock between CompiledFunction.lock and evalLock, and a data race on the compiler cache - #461

Merged
davidkoski merged 5 commits into
ml-explore:mainfrom
DanielMercado60660:compile-lock-fix-upstream
Aug 24, 2026
Merged

davidkoski merged 5 commits into
ml-explore:mainfrom
DanielMercado60660:compile-lock-fix-upstream

Conversation

@DanielMercado60660

Copy link
Copy Markdown
Contributor

Summary

Two threads calling compiled functions can deadlock permanently the first time a
trace runs. The per-instance CompiledFunction.lock and the process-global
evalLock are acquired in opposite orders on different threads. This makes
evalLock the outermost lock so there is only one order.

It is not a rare interleaving: with a cold compiler cache and one thread inside
compile_trace, any second thread entering a compiled function is enough. The
included reproduction wedges on round 1, every run, on both .cpu and .gpu.

While establishing that rule the PR also closes the one place that broke it:
CompiledFunction.deinit mutated the compiler cache without holding evalLock.
That is a data race rather than a deadlock, and it predates both commits below,
but it is the same cache and the same lock, so it is fixed here as its own
commit — see The compiler-cache erase in deinit.

The lock graph

CompiledFunction.call takes the per-instance lock, and then innerCall
takes the global evalLock and holds it through mlx_closure_apply. That gives
lock → evalLock. evalLock genuinely has to span that region: the
compiler cache it protects is a process-global singleton with no mutex
(mlx::core::detail::CompilerCache, compile.cpp), and the tracing flag is not
even per-thread — detail::InTracing::trace_stack() is a plain function-local
static, not thread_local (transforms.cpp), so while any thread traces,
in_tracing() is true for the whole process. Narrowing evalLock is not an
option.

But mlx_closure_apply on a cold cache entry runs compile_trace, which calls
back into the Swift body being compiled — and that body routinely calls other
compiled functions. Every MLXNN activation is a file-scope CompiledFunction,
so silu(x) inside any traced body enters compiledSilu.call and takes a second
per-instance lock while evalLock is held. That is evalLock → lock. The
two orders are inverted and close a cycle:

thread H (nested)                          thread W (plain)
  call(outer)                                call(shared)
    acquire L_cf(outer)                        acquire L_cf(shared)
    acquire evalLock
    mlx_closure_apply -> compile_trace
      -> back into Swift, body calls shared
    want L_cf(shared)  <-- held by W           want evalLock  <-- held by H

Note the per-instance lock is taken unconditionally, before C++ gets a chance
to short-circuit: detail::compile returns fun(inputs) immediately when any
input is_tracer(), so the nested call would not have touched the cache at all.
The Swift-side lock is the sole cause of the cycle.

This is first-use-only. compile_trace runs only when the cache entry is empty,
so once every trace is warm the window closes — and any new shape, dtype, stream
or constant set re-opens it.

Provenance

Neither commit is wrong on its own; the pair is.

Transforms+Vmap.swift has the same shape — evalLock held across a trace that
calls back into Swift — so it is on the evalLock → lock side too, and is
consistent with the ordering after this change.

Reproduction

New executable target CompileLockRepro (Source/CompileLockRepro). It is a
separate executable on purpose: a wedged round cannot be recovered and would take
a shared test process down with it.

It is deterministic rather than a race that has to be lost:

  • fresh compile() results per round, so the compiler cache is cold and the
    callback into Swift actually happens;
  • a semaphore signalled from inside the traced body, so the second thread is
    released only once the first is provably inside compile_trace holding
    evalLock;
  • a usleep in that same traced body holding the window open for 200 ms;
  • threads owned by the repro, so nothing depends on test-runner parallelism;
  • a 10 s per-round deadline, so a hang is a fast hard failure rather than a
    stuck CI job. Exit 1 means deadlock, exit 2 means the trace never ran (e.g.
    MLX_DISABLE_COMPILE), so a pass cannot be a false negative.
swift build --product CompileLockRepro
.build/debug/CompileLockRepro          # --device cpu|gpu

SwiftPM (command line) does not build the Metal shaders and the MLX scheduler
opens a GPU stream when constructed, so a swift build product needs an
mlx.metallib beside the binary even for --device cpu; one built by
xcodebuild -project xcode/MLX.xcodeproj -scheme MLX works.

Before:

compile-lock repro: device=Device(gpu, 0), 20 rounds, 10s deadline per round
FAIL: round 1: thread H did not complete within 10s.
  ...
  enteredTrace=true
  This is the lock-order inversion; the process is deadlocked.

sample(1) of the wedged process, 1572/1572 samples in identical frames:

Thread H-nested-compiled-call
  CompiledFunction.call                Transforms+Compile.swift:40   holds L_cf(outer)
  CompiledFunction.innerCall           Transforms+Compile.swift:112  holds evalLock (:89)
  mlx_closure_apply -> detail::compile compile.cpp:1180 -> :1126
  detail::compile_trace                compile.cpp:408
  [back into Swift] inner (tracers:)   Transforms+Compile.swift:70
  <traced body> -> CompiledFunction.call  Transforms+Compile.swift:40
  NSLocking.withLock -> __psynch_mutexwait

Thread W-plain-compiled-call
  CompiledFunction.call                Transforms+Compile.swift:40   holds L_cf(shared)
  CompiledFunction.innerCall           Transforms+Compile.swift:89
  -[NSRecursiveLock lock] -> __psynch_mutexwait

After: PASS: 20 rounds completed, no deadlock (exit 0), on .cpu and .gpu.

The fix

Four commits.

1. Take evalLock outermost in CompiledFunction.call and drop the
acquisition inside innerCall. One global order, evalLock → lock, so the
inversion cannot form. evalLock is recursive, so the nested acquisition from
inside a trace is free.

The cost is small: innerCall already acquired evalLock before doing any real
work and held it through mlx_closure_apply — that is, across the evaluation
itself — so this widens the held region only by the state gathering and closure
creation at the top of innerCall.

lock also becomes an NSRecursiveLock, so a compiled function whose body
re-enters that same compiled function on one thread no longer self-deadlocks on
a non-recursive NSLock.

2. Skip the per-instance lock when the function has no state. The
_updateInternal swaps in innerCall are the only mutation of anything
reachable from self — every other stored property is let and id is written
once at init — so with inputs and outputs both empty there is nothing for
lock to guard, and evalLock (held by commit 1) already serializes every
caller. This is the common case: every compiled function in MLXNN is stateless.
Optional; the deadlock is fixed without it. Happy to drop it if you would rather
keep the change minimal.

3. Check the ordering in debug builds. A passing repro only shows the race
did not fire. evalLock is recursive, so a successful try() cannot tell "this
thread owns it" from "nobody owns it"; ownership is recorded explicitly instead.
withEvalLock(_:) is now the single choke point for evalLock and keeps a
per-thread depth (pthread-specific, so the bookkeeping needs no lock of its own —
it runs while holding evalLock). CompiledFunction acquires its lock through
withInstanceLock, which requires that depth to be non-zero. All of it compiles
out of release builds.

Moving the per-instance lock back outside evalLock trips it on the first
compiled call, deterministically. Verified by doing exactly that locally:

MLX/Transforms+Eval.swift:66: Assertion failed: CompiledFunction.lock was
acquired without holding evalLock; that is the lock-order inversion, see
CompiledFunction.call

4. Erase the compiler cache under evalLock in deinit. See below.

The compiler-cache erase in deinit

CompiledFunction.deinit called mlx_detail_compile_erase(id) with no lock
held. That is compiler_cache().erase(fun_id) — a std::unordered_map::erase on
the same unguarded process-global cache that innerCall reads and inserts into
through CompilerCache::find, whose first statement is cache_[fun_id] (itself
an insert that can rehash) and whose slow path is entries.push_back. Worse,
find hands back a reference to an entry that the caller then fills in and
reads from across the entire trace and compile_replace. An unlocked erase can
rehash the map out from under a concurrent find, or destroy the very entry
another thread is compiling into.

This is a pre-existing data race — it is older than either commit in
Provenance — and it is not the deadlock. It is in this PR because
it is the same cache and the same lock: with commits 1–3 the rule is "every use
of the compiler cache happens under evalLock", and deinit was the one site
that broke it. Leaving it out would land a debug ordering checker that
deliberately does not check one of the two sites that mutate the cache.

It is commit 4 and touches nothing the first three touch beyond the same file, so
it can be dropped or split without disturbing them.

deinit runs on the releasing thread — what can that thread be holding?

Taking a lock in a deinit deserves the enumeration, since deinit fires
wherever the last reference happens to be dropped:

  • A thread holding no MLX lock. It waits for the eval or compiled call in
    flight and then erases. Bounded — evalLock is only ever held for the duration
    of one call — but it is new blocking; see the note on cost below.
  • A thread already inside a compiled call or an eval. A traced body that
    drops the last reference to some other compiled function is the realistic
    case. That thread already holds evalLock, which is an NSRecursiveLock, so
    the acquisition is free and cannot self-deadlock.
  • A thread holding some CompiledFunction's per-instance lock. After
    commit 1 that implies it holds evalLock too, so this is the previous case. No
    new edge is added to the lock graph.
  • A back-end worker thread, while an evalLock holder waits for it. This
    would be a genuine deadlock, and it does not arise: nothing in mlx-swift hands
    a CompiledFunction reference to one. The only strong reference is the closure
    returned by compile(...), and the only other retain is the
    ClosureCaptureState built by new_mlx_closure in innerCall, whose
    mlx_closure_free runs in that same function's defer, on the calling thread,
    with evalLock held. Nothing Swift-owned is stored in a CacheEntry — on the
    mlx_detail_compile path entry.extra is always nullptr — nor in the tape.
  • Process teardown. Unchanged in kind: a thread wedged holding evalLock
    already hangs atexit, which is why the repro uses _exit. MLXNN's
    file-scope compiled functions are Swift globals and are not destroyed at exit.

The residual hazard is on the caller's side, and it is worth stating: after this
change, releasing a compiled function is a call into MLX. Dropping one while
holding a lock that an evalLock holder needs is an inversion, exactly as
calling eval() there would be. That is the existing discipline for this library,
but a nil assignment advertises it much less loudly than eval() does.

Why the erase stays synchronous

The obvious way to avoid the blocking would be to hand the erase to a queue. That
is not safe here: id is the object's own address
(Unmanaged.passUnretained(self).toOpaque()), so once the object is freed a new
CompiledFunction can be allocated at the same address and receive the same id.
A deferred erase could delete the new function's cache entries. Doing it in
deinit, before the memory is released, is what keeps the id unambiguous.

Note also that detail::compile_erase has a second caller in compile.cpp — the
shared_ptr deleter in the public mlx::core::compile() — which is likewise
unlocked and can run on any thread. mlx-swift never reaches it (it calls only
mlx_detail_compile), so CompiledFunction.deinit is the only erase site this
package can fix from Swift. If you would rather have the lock inside
CompilerCache itself, that is a change on the C++ side and would cover both.

Tests

Tests/MLXTests/CompileLockOrderingTests.swift — none of these can hang, all are
safe in the shared test process:

  • the ownership bookkeeping can report "not held", and is per thread — otherwise
    a violation count of zero would prove nothing;

  • a stateful compiled function really reaches the checked acquisition (the
    counters also assert the check ran at all, so a pass is not vacuous);

  • the nested case, one compiled function called from inside another's trace;

  • the erase interleaving: one thread is parked inside compile_trace holding
    evalLock — it signals a semaphore from inside the traced body and then holds
    the window open — while a second thread drops the last reference to a warmed
    compiled function at that instant. The test asserts that the second thread
    observes the traced body to have finished by the time its erase returns,
    which is what "serialized against find" means, and that exactly one erase
    happened at that point. Reverting deinit to the unlocked form fails it on the
    intended assertion, deterministically:

    CompileLockOrderingTests.swift:252: error: testDeinitEraseSerializesAgainstATraceInFlight
    XCTAssertFalse failed - the compiler-cache erase completed while another thread
    was still inside compile_trace -- the erase is not serialized against find/insert
    
  • the erase under load: 200 deinit-erases against 200 cold compiles on two
    owned threads, covering the cache_[fun_id] insert-and-rehash window as well
    as tracing, asserting the results stay correct and that the erase path ran.

What that pair does and does not establish, plainly: mutual exclusion here is
structural, not statistical. Every find/insert is inside withEvalLock (from
call, the only caller of innerCall), and after commit 4 every erase is too;
the tests show the erase path really is reached with evalLock held under the
exact interleaving that used to race, and that adding the acquisition does not
wedge. The load test is a stress loop, so a green run of it alone is not proof
— which is why its assertions are the counters rather than the absence of a
crash.

No sanitizer run is included. TSan works on this toolchain, but swift test --sanitize=thread rebuilds the ~300 vendored MLX C++ translation units under it,
and at the vendored 0.31.1 it would report the pre-existing unsynchronized
globals in transforms.cpp (trace_stack, grad_counter) throughout, which
buries rather than sharpens the signal for this one line. That reasoning stops
holding after a bump to mlx 0.32.x, where those globals are thread_local: with
the background noise gone a TSan run becomes considerably more attractive, and
is the natural way to re-check this lane against a bumped core.

swift test: 541 → 547 tests, 0 failures before and after (541 existing on upstream main plus
the 6 new). CompileLockRepro: fails on round 1 before, 20/20 rounds after, on
.cpu and .gpu.

There were previously no concurrency tests in the package
(grep -rn concurrentPerform Tests/ Source/ is empty), which is presumably why
this went unnoticed.

Not fixed here

By design rather than by bug: because evalLock covers mlx_closure_apply,
every compiled call serializes against every other compiled call and every
eval/asyncEval in the process. This PR does not change that — it only moves
the acquisition a few statements earlier — but it does put a ceiling on
concurrent multi-stream use that is worth being explicit about. Commit 4 extends
the same serialization to the release of a compiled function.

The package vendors mlx 0.31.1, where the compiler cache and the tracing state
are process-global: CompilerCache has no mutex and
detail::InTracing::trace_stack() is a plain function-local static rather than
thread_local. evalLock is a Swift-side workaround for both, which is why it
has to span the whole trace rather than just the evaluation. That is no longer
the upstream shape. mlx 0.32.0 made the tracing stack, InTracing::grad_counter,
RetainGraph::tracing_counter and the compiler-cache singleton itself
thread_local, and 0.32.1 redesigned the cache — detail::CompileCache, an
unordered_map<uintptr_t, shared_ptr<vector<CacheEntry>>>, a find that returns
the entry reference together with an owning shared_ptr that keeps it alive
across the trace, and compile_erase(const weak_ptr<CompileCache>&, uintptr_t)
so a deleter erases from the cache of the thread that compiled, or from none at
all if that thread is gone. After a core bump to 0.32.x, evalLock would no
longer need to cover the trace for correctness and could shrink back to guarding
evaluation.

None of that supersedes this PR. The lock-order inversion fixed here is entirely
Swift-side — two Swift locks taken in opposite orders on different threads — and
it stands at any core version: a per-thread cache removes the C++ data race but
not the cycle. As context rather than a request, one cost of the thread-local
design is that nothing is reused across threads, so N worker threads pay N traces
and hold N tapes — which is why a shared cache under a mutex would still interest
embedders whose compiled functions are process-wide singletons, as every MLXNN
activation is.

Noted separately while reading the same package, and not fixed in this PR:
StreamOrDevice.stream(_:) (Source/MLX/Stream.swift:54-56) discards its
argument — public static func stream(_ stream: Stream) -> StreamOrDevice { StreamOrDevice(Device.defaultStream()) }
— so every caller that names a stream explicitly silently gets the device's
default stream instead. It is upstream-report shaped: a one-line fix
(StreamOrDevice(stream)), no relationship to the lock inversion, and its own
behavioural blast radius, since any code that has been relying on the observed
behaviour would start running where it asked to. Worth filing beside this PR
rather than inside it.

Related issue on ml-explore/mlx: ml-explore/mlx#4377

Daniel and others added 5 commits August 22, 2026 21:23
CompiledFunction.call takes the per-instance lock (ml-explore#226) and then, in
innerCall, the process-global evalLock (ml-explore#339), holding it across
mlx_closure_apply.  On a cold cache entry that runs compile_trace, which
calls back into the Swift body being compiled; if that body calls another
compiled function -- every MLXNN activation does, e.g. silu calls the
file-scope compiledSilu -- the nested call takes a second per-instance
lock while evalLock is held.

    thread H:  L_cf(outer)  -> evalLock     -> wants L_cf(shared)
    thread W:  L_cf(shared) -> wants evalLock

The orders are inverted and the two threads deadlock permanently.

CompileLockRepro owns its threads, uses fresh compile() results per round
so the compiler cache is cold, signals from inside the traced body so W is
ordered after H enters the trace, and holds the window open with a usleep
in that body.  A per-round deadline turns the hang into a hard failure.

It is its own executable because a wedged round cannot be recovered and
would take a shared test process down with it.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Establishes one global lock order, evalLock -> CompiledFunction.lock, so
the inversion described in the previous commit cannot form.

The reverse order is not avoidable on its own: the trace has to run under
evalLock (the compiler cache is an unguarded global and the tracing flag is
a plain static, not thread_local), the trace calls back into Swift, and the
traced body calls other compiled functions -- every MLXNN activation is a
file-scope CompiledFunction.  So evalLock -> lock happens regardless.  The
only fix available is to make it the *only* order, which is what taking
evalLock at the top of call(_:) does.

evalLock is recursive, so the nested acquisition from inside a trace is
free.  innerCall already took evalLock before doing any real work and held
it through mlx_closure_apply, so this widens the held region only by the
state gathering and closure creation at the top of innerCall.

lock becomes an NSRecursiveLock so that a compiled function whose body
re-enters that same compiled function on one thread no longer self-
deadlocks on a non-recursive NSLock.

CompileLockRepro: 20/20 rounds pass; without this it wedges on round 1.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
ml-explore#226 added CompiledFunction.lock to guard the observed state.  The
_updateInternal swaps in innerCall are the only mutation of anything
reachable from self -- every other stored property is let, and id is
written once at init -- so when inputs and outputs are both empty the lock
has nothing to guard.  With the previous commit evalLock is held for the
whole call, which already serializes the compiler cache and every other
caller, so dropping the acquisition changes nothing observable.

Optional on top of the previous commit; the deadlock is fixed without it.
It removes a mutex round-trip from the hot path of the common case: every
compiled function in MLXNN is stateless.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
A passing CompileLockRepro only shows that the race did not fire, so assert
the ordering itself.

evalLock is recursive, and a successful try() on a recursive lock cannot
distinguish 'this thread owns it' from 'nobody owns it', so ownership has to
be recorded explicitly.  withEvalLock is now the single choke point for
evalLock; in debug builds it keeps a per-thread depth (pthread-specific, so
the bookkeeping itself needs no lock -- it runs while holding evalLock).

CompiledFunction acquires its per-instance lock through withInstanceLock,
which requires that depth to be non-zero.  Moving the per-instance lock back
outside evalLock -- the shape that deadlocks -- trips it on the first
compiled call, with no race to lose.  All of it compiles out of release
builds.

CompileLockOrderingTests covers: that the bookkeeping can report 'not held'
and is per thread (otherwise a violation count of zero would prove nothing),
that a stateful compiled function really reaches the checked acquisition,
and the nested case where one compiled function is called from inside
another's trace.

Verified by inverting the order locally: the assertion fires immediately
with 'CompiledFunction.lock was acquired without holding evalLock'.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
`mlx_detail_compile_erase` is `compiler_cache().erase(fun_id)` on the
process-global `mlx::core::detail::CompilerCache`, which has no
synchronization of its own.  `CompiledFunction.deinit` called it with no
lock held, so it could run concurrently with another thread's
`CompilerCache::find` -- `cache_[fun_id]`, which can itself insert and
rehash, followed by `entries.push_back` -- and with the whole trace that
fills in the entry `find` returned a reference to.  Unlike the lock-order
inversion fixed in the preceding commits, this is a data race rather than
a deadlock, and it predates them.

Every other use of that cache already happens under `evalLock`, so the
erase takes it too.  `evalLock` is recursive, so a deinit that fires on a
thread already inside a compiled call or an `eval` -- a traced body
dropping the last reference to some other compiled function, for instance
-- re-enters for free, and a thread holding a `CompiledFunction.lock`
holds `evalLock` already, so no new edge is added to the lock graph.  A
deinit on a thread with no MLX lock simply waits for the call in flight.
Nothing in the library hands a `CompiledFunction` reference to a back-end
worker thread, so this erase can never be what an `evalLock` holder is
waiting for.

The erase stays synchronous rather than deferred: `id` is the object's own
address, so a later `CompiledFunction` can be allocated at the same
address and receive the same id, and a deferred erase could delete that
new function's entries.

Two tests.  The first is the interleaving itself: one thread is parked
inside `compile_trace` holding `evalLock` (it signals a semaphore from
inside the traced body), the other drops the last reference to a warmed
compiled function at that instant, and the erase is required to complete
only after the first thread has left the traced body -- which is what
serialization against `find` means.  The second runs a stream of erases
against a stream of cold compiles, covering the insert-and-rehash window
too, and checks the results are still correct.  Both assert the debug
counters so a pass cannot be vacuous.

`swift test`: 530 -> 532 tests, 0 failures.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
@davidkoski

Copy link
Copy Markdown
Member

Nice find!

Comment thread Package.swift
// Deterministic reproduction of the CompiledFunction/evalLock
// lock-order inversion. A wedged run deadlocks its threads
// permanently, so it is deliberately kept out of the test bundle and
// built as its own executable.

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I wonder if this could be done in swift testing with the subprocess capability? Perhaps an exercise for the future.

@davidkoski davidkoski left a comment

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Changes look good, thank you for finding this!

@davidkoski
davidkoski merged commit df9ae26 into ml-explore:main Aug 24, 2026
7 checks passed
jeethu added a commit to numen-tech/mlx-swift that referenced this pull request Aug 26, 2026
Brings the fork up to upstream main as of 2026-08-26: the eval deadlock fix (ml-explore#461),
the MLXArray(rawPointer:) finalizer leak fix (ml-explore#448), nested initializers (ml-explore#456),
CudaBuild plugin gated to Linux (ml-explore#447), CUDA build fixes (ml-explore#442/ml-explore#451) and the
update-mlx.sh tweak (ml-explore#445). The vendored core submodule stays on
numen-tech/mlx prism-0.31.1-fixes (b2b8a3d); no 0.32 mlx-c bindings exist yet.
bri-prism added a commit to PrismML-Eng/mlx-swift that referenced this pull request Sep 22, 2026
* fix(Device): resolve initial default from the MLX C++ core, not hard-coded GPU (ml-explore#422)

Closes ml-explore#395.

`_resolveGlobalDefaultDevice()` returned `.gpu` whenever no explicit
`setDefault(...)` had been called, so on a CPU-only host (Linux with no
Metal / no CUDA) `Device.defaultDevice()` and `StreamOrDevice.default`
both returned a device that simply doesn't exist, causing every
subsequent op routed to the default stream to fail.

The deprecated `Device.init()` already shows the right pattern: call
`mlx_get_default_device()` and let the C++ core pick the device its
build supports. Use that path instead of the hard-coded `.gpu`.

* Allowing SPM to compile on Linux with CUDA (ml-explore#413)

* Support CUDA compilation with SPM
* Enable CUDA by default on Linux
* Add code generation capabilities in CudaBuild plugin
* Use `tools/update-mlx.sh` to generate `cuda_jit_sources.h`

* Add MultiOptimizer (port from Python mlx.optimizers) (ml-explore#431)

Port `mlx.optimizers.MultiOptimizer` to Swift: a meta-optimizer that routes
each model parameter to the first sub-optimizer whose filter predicate matches
its flattened key path, with the last optimizer acting as the unconditioned
fallback. This allows using different optimizers or hyperparameters for
different subsets of a model's weights.

`update(model:gradients:)` partitions the flattened gradient tree by
first-matching filter and delegates each subset to the corresponding child
optimizer; `innerState()` surfaces the children's state. It conforms to the
`Optimizer` protocol directly since it holds no per-parameter state of its own.

Closes ml-explore#228

Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* Add MLXArray and DType extensions (ml-explore#429)

* Add MLXArray and DType extensions
* Simplify magic numbers
* Return MLXArray instead of Double

* Add Muon optimizer (port from Python mlx.optimizers) (ml-explore#432)

Port `mlx.optimizers.Muon` (MomentUm Orthogonalized by Newton-Schulz): a
momentum optimizer that orthogonalizes the update for >=2D parameters via a
5-step quintic Newton-Schulz iteration, scaling the learning rate by
sqrt(max(1, rows/cols)). 0D/1D parameters fall back to a plain momentum/Nesterov
update. Mirrors the reference (weight decay, Nesterov default on, ns_steps=5).

Tests pin the 1D momentum/Nesterov math exactly and exercise the 2D
orthogonalization path (shape preserved, finite, parameter moved).

Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* Add learning-rate schedules (port from Python mlx.optimizers.schedulers) (ml-explore#433)

Port the mlx.optimizers schedule functions to Swift: exponentialDecay,
stepDecay, cosineDecay, linearSchedule, and joinSchedules. Each returns a
(Int) -> Float mapping a step to a learning rate, composing with the mutable
`learningRate` exposed by the optimizers (assign optimizer.learningRate =
schedule(step) each step). mlx-swift previously had no LR schedules.

Tests pin exact values against the reference docstring examples.

Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

* Move .complex64 to .float32 (ml-explore#434)

* guard use of Process (ml-explore#437)

- does not build on iOS, etc.
- fix ml-explore#436

* Add fp8 conversion utilities (ml-explore#435)

* Add FP8 conversion utilities.

* build iOS in CI (ml-explore#440)

- add an iOS build in CI to catch iOS-only build breaks
- move the docs build into the same macOS runner

* Add global scale support to quantized layers (ml-explore#426)

* Add global scale support to quantized layers

Store optional globalScale on QuantizedLinear and QuantizedEmbedding so nvfp4 weights can preserve the scale needed by lower-level MLX quantize/dequantize operations.

Forward the scale when creating or dequantizing weights, add a direct pre-quantized QuantizedEmbedding initializer to match QuantizedLinear, and guard global-scale execution paths on Metal because MLX does not support globalScale dequantization there.

Add focused tests for the new layer state and parameter exposure without broadening the generic quantization API surface.

* Keep nvfp4 global scale on the quantized path for Metal

Since MLX encodes nvfp4 group scales pre-multiplied by (448 * 6) / globalScale,
the global scale factors out of dequantization linearly. Exploit this by always
running quantizedMM (supported on Metal without globalScale) and applying
globalScale / (448 * 6) to the output, instead of dequantizing the full weight
and falling back to a plain matmul.

This keeps weights quantized end-to-end on all backends, removes the
dequantizedWeight fallback from QuantizedLinear and QuantizedEmbedding, and
applies the scale before the bias add so the bias is unaffected.

Add Metal tests that validate the GPU quantized path against CPU
dequantized(globalScale:) references for QuantizedLinear (matmul + bias) and
QuantizedEmbedding (asLinear and index lookup).

* Fix FinalizerCaptureState leak in MLXArray(rawPointer:_:dtype:finalizer:) (ml-explore#448)

The initializer retains a box holding the finalizer closure and hands it to
mlx as an opaque payload, but `finalizerTrampoline` read it back with
`takeUnretainedValue`, which never consumes the `passRetained` +1. The local
`free(ptr:)` that would have balanced it was never passed anywhere — only the
trampoline is — so it was dead code and the box leaked.

The box owns the closure, so the closure's captures leak with it. The
initializer's own documented example captures an `IOSurface` and comments that
it "implicitly releases it when it returns"; it did not, and the surface was
pinned for the lifetime of the process. `testIOSurface` covers the path but
only prints, so nothing caught it.

`takeRetainedValue` consumes the +1 exactly once. mlx-c documents the dtor as
"Callback for when the buffer is no longer needed" and calls it once per
buffer, and the trampoline has exactly one caller, so a single consume is
correct. Removes the dead `free(ptr:)`.

Adds a differential regression test: with the bug the finalizer runs but its
captures are never released (witness releases 0); with the fix they are (1).
The test drains mlx's allocator cache first, since a freed buffer's dtor is
not guaranteed to have run at scope exit, and uses lock-guarded counters
because mlx may invoke the dtor from its own thread.

Full suite: 534 tests, 0 failures.

Co-authored-by: Fred Jura <digital@afya.ai>

* Derive the jit-source lists in update-mlx.sh instead of hand-listing them (ml-explore#445)

`tools/update-mlx.sh` hand-tracked two lists that mlx's own CMake derives, so
each one silently went stale whenever the vendored mlx submodule gained a jit
source.

The CUDA `MLX_JIT_SOURCES` argument named 16 `device/` headers explicitly, where
`mlx/backend/cuda/CMakeLists.txt` uses `file(GLOB)` over `device/*.h` and
`device/*.cuh`. A bump that adds one regenerated a `cuda_jit_sources.h` missing
its array, failing later as `use of undeclared identifier 'jit_source_*'`:
mlx v0.31.2 needs 17 and v0.32.0 needs 22, against the 16 the frozen list
produced.

The Metal step named 46 `make` targets explicitly, where
`mlx/backend/metal/CMakeLists.txt` declares them via `make_jit_source()`. A
bump that adds one left `jit/includes.h` declaring a
`mlx::core::metal::<name>()` that no generated source defined, i.e. an
undefined symbol at link time on macOS and iOS: v0.31.2 declares 47 and
v0.32.0 declares 48, against the same 46.

Both lists are now derived from the vendored tree at run time. Against the
currently pinned mlx (v0.31.1) the derived values are identical to the removed
literals, and the regenerated `cuda_jit_sources.h` is byte-for-byte identical
to the committed one, so this is a no-op at the current pin.

Each derivation is guarded: with `set -e` an empty result would not fail on its
own, and bare `make` with no targets would quietly build `all` instead.

* Improve CUDA NVCC spawn performance (on Linux) (ml-explore#451)

* Reap encuda's child processes with waitpid instead of polling isRunning

`waitUntilExitWorkaround()` polled `Process.isRunning` in a 50ms sleep
loop, added because `waitUntilExit()` "never returns" on aarch64 Linux.
It did not fix that hang -- `isRunning` reads the same Foundation
bookkeeping that stalls, so the workaround only changed a blocked call
into an infinite sleep loop.

Observed during a parallel SwiftPM build of the CUDA backend for
linux/arm64 (aarch64-unknown-linux-gnu, CUDA 13, sm_121): `nvcc` finishes
and writes its output in full, then becomes a zombie that is never reaped,
while `encuda-tool` spins in `hrtimer_nanosleep` holding a completed 15MB
transpiled .cpp on disk. Invoking encuda directly on the same file with
identical arguments completes in 2-3 seconds, so the work itself is fine
-- the process simply never notices the child finished.

It is load-dependent: `--jobs 2` transpiled all 96 .cu files without
incident, while an unconstrained build (SwiftPM defaults to every core)
deadlocked. That is consistent with pipe file descriptors leaking between
concurrently spawned siblings -- the hung `encuda-tool` held both the read
and the write end of the same pipes, so no reader could ever see EOF.

`waitpid` answers the only question this code has: did the child exit, and
with what status. It does not depend on Foundation's reaper, and it
collects the zombie. On ECHILD -- Foundation's own reaper winning the race
-- it falls back to `terminationStatus`, which is authoritative in exactly
that case.

Exit status is decoded without the C macros, which Swift does not import:
the low 7 bits carry the terminating signal and the next 8 the exit code,
an encoding shared by Linux and Darwin. Signals are reported as
`128 + signal`, following the shell convention, so a killed compiler is
still a non-zero failure rather than a silent success.

Verified on macOS: exit codes 0, 7 and 42 round-trip exactly; SIGKILL and
SIGTERM report 137 and 143; no reapable child remains afterwards. The
aarch64 Linux deadlock itself has not been reproduced locally -- it needs
a CUDA toolchain and a parallel build to trigger -- so this is a fix for a
diagnosed root cause rather than one observed to clear the symptom.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01K1rzxACJqmxVqgRsa4xj3R

* Don't read terminationStatus before Foundation has published it

The waitpid change cleared the deadlock it was written for -- an
unconstrained SwiftPM build of mlx's CUDA backend for aarch64 Linux
(CUDA 13.2, sm_121, 16 jobs) now runs through the transpile step at
full speed instead of stalling forever on a zombie nvcc. But it
replaced the hang with a crash.

`waitpid` and Foundation's own reaper thread race for every child.
When Foundation wins, our waitpid returns ECHILD and the fallback read
`terminationStatus` immediately. That is not safe: Foundation
publishes the result in two steps, and reading the status inside the
window before `isRunning` goes false traps the process.

    *** Program crashed: System trap at 0x0000ffff90b6fb50 ***
    Thread 0 "encuda-tool" crashed:
      0  Process.terminationStatus.getter + 80 in libFoundation.so
      1  Process.waitForExitStatus() + 107 in encuda-tool
           at Source/Encuda/encuda-utils.swift:73:24
      2  Encuda.Compile.run() + 1883 in encuda-tool

Roughly one transpile in 250 died this way, at a different .cu file
each run. The build failure it produces is unusually opaque: the crash
report goes to the transpiler's own output, so `swift build` exits 1
with no diagnostic at all, and the last thing in the log is an
ordinary "Compiling scan.cu to scan.cpp" progress line.

So wait for `isRunning` to clear before reading the status -- but on a
deadline, since an unbounded wait is the hang this function exists to
avoid. On expiry it reports 254 rather than trapping or spinning: the
caller turns any non-zero status into a build error, which is
diagnosable, and neither a trap nor a hang is.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01RdegWrNm5PKBqeXdkgmn9Q

---------

Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>

* Condition CUDA build plugin on Linux host (ml-explore#447)

* Nested initializers  (ml-explore#456)

* nested initializers

* [CUDA] Build the CUDA argument lists in steps (ml-explore#442)

Three argument lists in the CUDA build are assembled as long `+` chains that the
type checker struggles with. On a Swift 6.2 toolchain they exceed its budget
outright:

    encuda-compile.swift:66:17: error: the compiler is unable to type-check this
    expression in reasonable time; try breaking up the expression into distinct
    sub-expressions
    plugin.swift:172:13: error: the compiler is unable to type-check this
    expression in reasonable time; ...

Appends into local arrays instead, which is behaviourally identical and is what
the diagnostic suggests. The plugin's two chains move into `compileArguments` and
`linkArguments` helpers to keep `createBuildCommands` readable, and adding a
further argument to either list no longer risks tipping the type checker over.

Note this does *not* make the package buildable on Swift 6.2: SwiftPM 6.2 ignores
non-Swift files produced by a build-tool plugin ("Only Swift is supported for
generated plugin source files at this time"), so every generated .cpp is dropped
and the CUDA symbols go undefined at link time. The 6.3 floor is load-bearing.
This is purely a type-check-cost and readability change.

Co-authored-by: Claude Opus 5 <noreply@anthropic.com>

* Fix a deadlock between `CompiledFunction.lock` and `evalLock`, and a data race on the compiler cache (ml-explore#461)

* add a deterministic repro for the compile/eval lock-order inversion

CompiledFunction.call takes the per-instance lock (ml-explore#226) and then, in
innerCall, the process-global evalLock (ml-explore#339), holding it across
mlx_closure_apply.  On a cold cache entry that runs compile_trace, which
calls back into the Swift body being compiled; if that body calls another
compiled function -- every MLXNN activation does, e.g. silu calls the
file-scope compiledSilu -- the nested call takes a second per-instance
lock while evalLock is held.

    thread H:  L_cf(outer)  -> evalLock     -> wants L_cf(shared)
    thread W:  L_cf(shared) -> wants evalLock

The orders are inverted and the two threads deadlock permanently.

CompileLockRepro owns its threads, uses fresh compile() results per round
so the compiler cache is cold, signals from inside the traced body so W is
ordered after H enters the trace, and holds the window open with a usleep
in that body.  A per-round deadline turns the hang into a hard failure.

It is its own executable because a wedged round cannot be recovered and
would take a shared test process down with it.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

* take evalLock outermost in CompiledFunction.call

Establishes one global lock order, evalLock -> CompiledFunction.lock, so
the inversion described in the previous commit cannot form.

The reverse order is not avoidable on its own: the trace has to run under
evalLock (the compiler cache is an unguarded global and the tracing flag is
a plain static, not thread_local), the trace calls back into Swift, and the
traced body calls other compiled functions -- every MLXNN activation is a
file-scope CompiledFunction.  So evalLock -> lock happens regardless.  The
only fix available is to make it the *only* order, which is what taking
evalLock at the top of call(_:) does.

evalLock is recursive, so the nested acquisition from inside a trace is
free.  innerCall already took evalLock before doing any real work and held
it through mlx_closure_apply, so this widens the held region only by the
state gathering and closure creation at the top of innerCall.

lock becomes an NSRecursiveLock so that a compiled function whose body
re-enters that same compiled function on one thread no longer self-
deadlocks on a non-recursive NSLock.

CompileLockRepro: 20/20 rounds pass; without this it wedges on round 1.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

* skip the per-instance lock when a compiled function has no state

ml-explore#226 added CompiledFunction.lock to guard the observed state.  The
_updateInternal swaps in innerCall are the only mutation of anything
reachable from self -- every other stored property is let, and id is
written once at init -- so when inputs and outputs are both empty the lock
has nothing to guard.  With the previous commit evalLock is held for the
whole call, which already serializes the compiler cache and every other
caller, so dropping the acquisition changes nothing observable.

Optional on top of the previous commit; the deadlock is fixed without it.
It removes a mutex round-trip from the hot path of the common case: every
compiled function in MLXNN is stateless.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

* check the evalLock -> CompiledFunction.lock ordering in debug builds

A passing CompileLockRepro only shows that the race did not fire, so assert
the ordering itself.

evalLock is recursive, and a successful try() on a recursive lock cannot
distinguish 'this thread owns it' from 'nobody owns it', so ownership has to
be recorded explicitly.  withEvalLock is now the single choke point for
evalLock; in debug builds it keeps a per-thread depth (pthread-specific, so
the bookkeeping itself needs no lock -- it runs while holding evalLock).

CompiledFunction acquires its per-instance lock through withInstanceLock,
which requires that depth to be non-zero.  Moving the per-instance lock back
outside evalLock -- the shape that deadlocks -- trips it on the first
compiled call, with no race to lose.  All of it compiles out of release
builds.

CompileLockOrderingTests covers: that the bookkeeping can report 'not held'
and is per thread (otherwise a violation count of zero would prove nothing),
that a stateful compiled function really reaches the checked acquisition,
and the nested case where one compiled function is called from inside
another's trace.

Verified by inverting the order locally: the assertion fires immediately
with 'CompiledFunction.lock was acquired without holding evalLock'.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

* erase the compiler cache under evalLock in CompiledFunction.deinit

`mlx_detail_compile_erase` is `compiler_cache().erase(fun_id)` on the
process-global `mlx::core::detail::CompilerCache`, which has no
synchronization of its own.  `CompiledFunction.deinit` called it with no
lock held, so it could run concurrently with another thread's
`CompilerCache::find` -- `cache_[fun_id]`, which can itself insert and
rehash, followed by `entries.push_back` -- and with the whole trace that
fills in the entry `find` returned a reference to.  Unlike the lock-order
inversion fixed in the preceding commits, this is a data race rather than
a deadlock, and it predates them.

Every other use of that cache already happens under `evalLock`, so the
erase takes it too.  `evalLock` is recursive, so a deinit that fires on a
thread already inside a compiled call or an `eval` -- a traced body
dropping the last reference to some other compiled function, for instance
-- re-enters for free, and a thread holding a `CompiledFunction.lock`
holds `evalLock` already, so no new edge is added to the lock graph.  A
deinit on a thread with no MLX lock simply waits for the call in flight.
Nothing in the library hands a `CompiledFunction` reference to a back-end
worker thread, so this erase can never be what an `evalLock` holder is
waiting for.

The erase stays synchronous rather than deferred: `id` is the object's own
address, so a later `CompiledFunction` can be allocated at the same
address and receive the same id, and a deferred erase could delete that
new function's entries.

Two tests.  The first is the interleaving itself: one thread is parked
inside `compile_trace` holding `evalLock` (it signals a semaphore from
inside the traced body), the other drops the last reference to a warmed
compiled function at that instant, and the erase is required to complete
only after the first thread has left the traced body -- which is what
serialization against `find` means.  The second runs a stream of erases
against a stream of cold compiles, covering the insert-and-rehash window
too, and checks the results are still correct.  Both assert the debug
counters so a pass cannot be vacuous.

* Fix StreamOrDevice.stream(_:) ignoring its argument (ml-explore#463)

StreamOrDevice.stream(_:) discarded the stream it was given and returned
the default stream of the default device instead, so operations that
passed an explicit stream silently ran somewhere else. Return the given
stream and add a regression test.

* test Module parameters and compile (ml-explore#465)

* test Module parameters and compile

* update for mlx v0.32.2 (ml-explore#450)

* update for mlx v0.32.0

mlx v0.32.2

- carry `Source/Cmlx/mlx.patch` for mlx's `Device::operator<` /
  `Stream::operator<`, which are not strict weak orderings and now run on
  every `eval()` via `std::set<Stream>`. libc++ aborts on them under the
  Xcode 27 SDK. Still broken upstream in v0.32.0 and on main, no issue filed.
  Applied by `tools/patch-mlx.sh` from `update-mlx.sh` and CI, since
  `git submodule update` restores a pristine checkout.

- streams became thread affine in v0.31.2, which broke the multi-threaded
  tests. Adopt mlx's `new_thread_unsafe_stream` (added upstream for language
  bindings that lock themselves -- mlx-swift serializes on `evalLock`) via a
  small C shim, as mlx-c does not bind it.

- FFT ops gained a `norm` parameter: add `FFTNorm` and thread it through,
  plus `fftfreq`/`rfftfreq`.

- wrap the new reducing slice-updates in `ArrayAtIndices` and add
  `exportToDot`/`graphDescription`/`NodeNamer`.

- fix `SEEK_END` in the in-memory IO stream -- it computed from the read
  position rather than the end of the buffer, which v0.32.0's
  `load_safetensors` size check exposed.

- route `cpu/jit_compiler.cpp` through `mlx-conditional`: it shells out via
  `std::system()`, unavailable on iOS.

- `gemv` is JIT compiled as of v0.32.0, so drop `gemv.metal` from the
  metallib in `fix-metal-includes.sh`. The jit source lists themselves no
  longer need touching -- ml-explore#445 derives them from the vendored tree, and the
  derived sets match what a hand-bump would have produced here (48 Metal
  targets, 22 CUDA sources).

* disable cuda runs for now -- need clear_streams exposed in mlx-c

---------

Co-authored-by: David Koski <dkoski@apple.com>

* Fix wired-memory ticket lifecycle during cancellation (ml-explore#471)

* Fix condition typo (ml-explore#473)

* fix(ops): fix convolve padRight for even-sized kernels in same mode (ml-explore#466) (ml-explore#467)

* Add countNonzero to the Swift API (ml-explore#479)

* Add countNonzero + docs
* Add tests

* Distributed support (ml-explore#482)

Add distributed communication and sharded neural network layers

Port mlx.core.distributed and distributed layer primitives to Swift:

- Backend & Primitives:
  - Enable the ring distributed backend (CPU socket-based communication).
  - Add MLXDistributed.Group with rank/size/split and 8 communication
    primitives (allSum, allGather, send, recv, etc.).
  - Default distributed operations to the CPU stream.
  - Make initialize and split throw on failure.

- Layers & Optimization:
  - Add AllToShardedLinear, ShardedToAllLinear, and quantized counterparts 
    (subclassing Linear / QuantizedLinear for seamless model replacement).
  - Implement shardLinear, shardInPlace (supporting predicate sharding), 
    sumGradients, averageGradients, and clipGradNormSharded.
  - Ensure correct boundary checks and integer scaling for quantized layers.

- Testing & Infrastructure:
  - Add a reusable multi-process test harness (DistributedHarness) for 
    spawning local ranks via loopback sockets.
  - Port Python ring tests (collectives, dtypes, donation, send/recv).
  - Port layer, predicate sharding, and gradient tests; enable in CI.

* add thin logging shim (ml-explore#484)

* add thin logging shim

* patch for Cuda builds (ml-explore#478)

* patch for Cuda builds

- local apply of ml-explore/mlx#4480 + other shutdown items from ml-explore/mlx#4506
- fixes crash on exit for cuda builds when using cmake

* pool Streams, fix Device inheritance (ml-explore#472)

* pool Streams, fix Device inheritance

- see ml-explore/mlx#2118
- Streams are finite and seem to leak OS/Metal resources
- implements a pool to reuse Streams -- withNewDefaultStream is fully usable

- much more care is taken with GPU devices
- using index != 0 now works in all cases
- and inherits properly when creating new streams

Co-authored-by: Nick Henderson <525247+nwh@users.noreply.github.com>

* replace and improve integration tests (ml-explore#477)

* replace and improve integration tests

- previously we used a python script to generate swift code and check output against known values from mlx (python)
- this is replaced with a new and more powerful script (claude rewrite)
- much more integration test coverage and tools to examine that

- fixes per problems found by new integration tests
- linspace output type, add dtype override
- linspace default output type is float32 (matches python)
- convolve .same padding side matches python (padRight)
- nanToNum default for +/- infinity honors optionals (mlx defaults)
- tensorDot axes default 1 -> 2 (matches python)

- MLXNN Pooling padWidths matches python
- MLXNN Alibi slop function matches python
- MLXNN GRU apply bias to match python
- MLXNN Lion optimizer default beta matches python
- MLXNN Adafactor moving average matches python
- MLXNN Muon zero power newton shulz matches python

* fix linspace disambiguation

* split integration tests into their own test target -- we need to run them with TF32 disabled so we can guarantee consistent results on NAX vs non NAX hardware

Co-authored-by: Nick Henderson <525247+nwh@users.noreply.github.com>

* Adapt data stream callbacks to the updated C binding contract

---------

Co-authored-by: tsushanth <78000697+tsushanth@users.noreply.github.com>
Co-authored-by: Gabriele Mondada <gab@ijk.ch>
Co-authored-by: GoodOlClint <goodolclint@gmail.com>
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Co-authored-by: Ronald Mannak <ronaldmannak@me.com>
Co-authored-by: David Koski <46639364+davidkoski@users.noreply.github.com>
Co-authored-by: Lucas Newman <lucasnewman@me.com>
Co-authored-by: Alessio Pollero <alessio.pollero@gmail.com>
Co-authored-by: agerjura <136923563+agerjura@users.noreply.github.com>
Co-authored-by: Fred Jura <digital@afya.ai>
Co-authored-by: Joannis Orlandos <joannis@orlandos.nl>
Co-authored-by: Dr. Thorsten Fuchs <thorfuchs@outlook.com>
Co-authored-by: Valentin Roussellet <vroussellet@apple.com>
Co-authored-by: DanielMercado60660 <153678010+DanielMercado60660@users.noreply.github.com>
Co-authored-by: Alessio Pollero <alessio.pollero@ventureone.ae>
Co-authored-by: Charlie Le <charlie_le@apple.com>
Co-authored-by: David Koski <dkoski@apple.com>
Co-authored-by: Lukas Zima <76228598+Silenterc@users.noreply.github.com>
Co-authored-by: Shubhransh Gupta <54713516+shubhransh-gupta@users.noreply.github.com>
Co-authored-by: Nick Henderson <525247+nwh@users.noreply.github.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants