Skip to content

update for mlx v0.32.2 - #450

Merged
davidkoski merged 8 commits into
ml-explore:mainfrom
CharlieTLe:bump-mlx-0.32.0
Sep 1, 2026
Merged

davidkoski merged 8 commits into
ml-explore:mainfrom
CharlieTLe:bump-mlx-0.32.0

Conversation

@CharlieTLe

@CharlieTLe CharlieTLe commented Aug 8, 2026 •

Copy link
Copy Markdown
Contributor

Proposed changes

Updates the vendored mlx to v0.32.0 and mlx-c to fba4470 (its latest; those bindings were generated against 0.31.2, so 0.32.0's new ops aren't exposed yet). MLX_VERSION is bumped in Package.swift and xcode/xcconfig/Cmlx.xcconfig.

Adaptations the bump requires:

  • FFT gained a norm parameter. Every mlx_fft_* signature changed, so this adds a public FFTNorm enum (.backward default, preserving current behaviour) and threads it through Source/MLX/FFT.swift. Also wraps the new fftfreq/rfftfreq.

  • SEEK_END in the in-memory IO stream was wrong — it computed from the read position rather than the end of the buffer. Latent until v0.32.0's load_safetensors added a size check via seek(0, end) + tell(), which made saveToData/loadArrays round-trips fail. SaveTests covers it.

  • cpu/jit_compiler.cpp now shells out via std::system(), which is unavailable on iOS. Routed through mlx-conditional the same way compiled.cpp already is, and excluded from the direct build so it isn't compiled twice.

  • gemv is JIT compiled as of v0.32.0 (jit_kernels.cpp gained get_gemv_kernel using metal::gemv()), so gemv.metal no longer needs compiling into the metallib and comes out of fix-metal-includes.sh. The jit source lists themselves need no change — Derive the jit-source lists in update-mlx.sh instead of hand-listing them #445 derives them, and the derived sets match what a hand bump would produce (48 Metal targets, 22 CUDA sources).

  • Package.swift excludes and xcodeproj target membership updated for added/removed upstream files (jaccl moved under lib/, nccl_stub gone, new cuda/dirs.cpp and qmm sources, new backend/common/metal_kernel.{h,cpp}), plus regenerated include/, include-framework/ and mlx-generated/.

New API now reachable from mlx-c fba4470:

  • the reducing slice-updates, wired into ArrayAtIndices so pure-slice at[] updates use mlx_slice_update_add/prod/min/max instead of a scatter (mirroring upstream's mlx_add_item), falling back to scatter otherwise
  • exportToDot / graphDescription / NodeNamer in a new Source/MLX/GraphUtils.swift

CMakeLists.txt declares mlx ahead of mlx-c so both build paths use the same mlx — mlx-c currently pins v0.31.2. That can be dropped once mlx-c pins 0.32.0 itself, e.g. via ml-explore/mlx-c#123.

Per review, the mlx.patch scaffolding and the local new_thread_unsafe_stream shim were both removed in 2d27240 — see the two threads below. tools/update-mlx.sh, README.md and MAINTENANCE.md are byte-identical to main again.

Known blockers

Two failures remain, both from upstream changes rather than from this change. Neither is fixed here.

  1. The ordering abort. Device::operator< and Stream::operator< are not strict weak orderings, and 0.31.2 started running them on every eval() via std::set<Stream> in eval_impl. libc++ validates the comparator under the Xcode 27 SDK and aborts wherever cpu and gpu streams are both open, so the test run aborts rather than reporting failures. Reported as ml-explore/mlx#4083 and fixed by #4086, but that is on main and unreleased — v0.32.0 is still the newest tag. This needs an mlx release containing #4086.

  2. Thread-affine streams. Streams became thread affine in 0.31.2 (#3348, #3281) and upstream treats it as intended (#3727 was closed pointing at a consumer-side fix in mlx-lm). Because Stream.gpu/.cpu are cached process-wide, MLXRandomTests/testRandomThreadsDifferent, testRandomThreadsSame and TransformTests/testVmapThreadSafety fail with There is no Stream(gpu, 0) in current thread. They pass on 0.31.1. Making mlx-swift's default streams resolve per thread would fix them and needs no new API, but the stream model wants its own design rather than being settled inside a version bump, so it's deferred to follow-up work — see the Stream.swift thread. It doesn't block this bump.

Verification

On macOS 27.0 (Apple Silicon), Xcode 27.0:

swift build succeeds
xcodebuild -scheme mlx-swift-Package -destination 'generic/platform=iOS' succeeds
xcode/MLX.xcodeproj (-scheme MLX, macOS) succeeds
cmake .. -G Ninja && ninja succeeds, libs and examples link, ./tutorial runs
pre-commit run --all-files clean
swift test aborts — blocker 1 above
scripts/verify-docs.sh fails — see below

With the patch still applied locally, the full suite was 554 tests / 0 failures, so the two blockers above are the only outstanding failures.

verify-docs.sh fails in ExtractAPI, which can't find any C++ stdlib header (algorithm, cmath, iostream) while parsing the Cmlx target. This is not from this change: it fails identically with my new headers removed, and it failed the same way before the bump on unmodified metal-cpp. It looks like Xcode 27 parsing C++ headers in C mode. I couldn't get a clean baseline run to prove it — SwiftPM rejected the manifest in a baseline worktree under the docc plugin — so worth confirming on CI's Xcode.

Checklist

Put an x in the boxes that apply.

  • I have read the CONTRIBUTING document
  • I have run pre-commit run --all-files to format my code / installed pre-commit prior to committing changes
  • I have added tests that prove my fix is effective or that my feature works
  • I have updated the necessary documentation (if needed)

On tests: FFTTests (norm scaling per mode, round-trips, fftfreq/rfftfreq), GraphUtilsTests, slice-update cases in ArrayAtTests including a slice-vs-scatter agreement check, and a SaveTests regression for the SEEK_END fix. On docs: the new public API carries DocC comments; no article changes are needed now that the stream work is out.

@davidkoski

Copy link
Copy Markdown
Member

There is an incoming 0.32.0 mlx-c in ml-explore/mlx-c#123. #445 should take care of the JIT issue.

Comment thread .github/workflows/pull_request.yml Outdated

- name: Patch mlx submodule
shell: sh
run: ./tools/patch-mlx.sh

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.

We should not need this -- we should rely on mlx and mlx-c tags. Anything that needs special handling should be done statically via tools at the time when we make the new release.

In particular SwiftPM can't apply patches (or make any source mods) so this won't really work.

@CharlieTLe CharlieTLe Aug 11, 2026 •

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

You're right, and the SwiftPM point settles it — removed in 2d27240. I'd wired the patch into the places this repo controls (update-mlx.sh, both workflows, the CMake PATCH_COMMAND, a README step), but a package consumer gets the submodule pristine and never runs any of it, so they'd build with the bug regardless. That's the path that actually matters and I hadn't covered it. There's no static option either: the patch touches headers inside the submodule, so there's nothing checked in for tools/ to bake in at release time. update-mlx.sh, README.md and MAINTENANCE.md are byte-identical to main again.

What that leaves: ml-explore/mlx#4086 (the ordering fix, from #4083) is on main but unreleased — v0.32.0 is still the newest tag and git tag --contains for it is empty. So on the Xcode 27 SDK, libc++ validates the comparator and eval() aborts out of std::set<Stream> in eval_impl. It isn't confined to a couple of tests; it fires wherever cpu and gpu streams are both open, so the run aborts rather than reporting failures. Builds are all fine — swift build, iOS, and the xcodeproj all succeed.

So this needs an mlx release containing #4086 before it can go green. Happy to hold it until then and re-pin, or leave it open as-is if you'd rather have the diff visible in the meantime.

Separately, I rebased onto #445 and dropped my hand-added jit-source entries — I checked the derived sets match what the hand bump produced (48 Metal targets, 22 CUDA sources), so the committed mlx-generated output is unchanged. One leftover I'd flag: CMakeLists.txt still declares mlx ahead of mlx-c to pin v0.32.0, because mlx-c currently pins v0.31.2 and the two build paths would otherwise disagree. That comes out once mlx-c pins 0.32.0 itself, e.g. via #123.

Comment thread Source/MLX/Stream.swift Outdated

public static let gpu = Stream(mlx_default_gpu_stream_new())
public static let cpu = Stream(mlx_default_cpu_stream_new())
/// Make a stream on `device` that can be used from any thread.

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.

We have been pretty loose about streams until now. I think we may need to:

  • manage a per task/thread stream, e.g. if you create a new task you get an implicit new stream unless you use some kind of withNewStream the of API
  • because of [BUG] crash while using new_stream mlx#2118 I think we may need to manage a pool of streams -- we can't create and destroy them at will

@CharlieTLe CharlieTLe Aug 11, 2026 •

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

Understood — reading this as direction rather than something to solve here. Leaving it out of this PR.

Recording where I got to, in case it's useful later. The minimal step looks like it needs no new API: mlx_default_gpu_stream_new() already calls default_stream(DeviceType::gpu), and default_stream_storage is static thread_local as of 0.31.2, so each call already returns the calling thread's stream. The only reason it breaks is that Stream.gpu/.cpu are static let, so the first thread to touch them wins. Making those (and StreamOrDevice.cpu/.gpu, and Device.defaultStream, which is a stored property) resolve per thread would fix the three failing tests without a shim or an mlx-c binding.

It wouldn't be a complete answer, which I think is your point: withNewDefaultStream installs one new_stream as a @TaskLocal, so a task that hops threads across an await still breaks, and it does nothing for #2118 — Stream(_ device:) still creates a stream per call. Pooling plus per-task streams is the real fix.

So for this PR the three thread tests fail (MLXRandomTests/testRandomThreadsDifferent, testRandomThreadsSame, TransformTests/testVmapThreadSafety — all pass on 0.31.1). I've left them failing rather than skipping them, so the behaviour change stays visible; happy to add XCTSkips pointing at this thread if you'd rather.

One correction to what I wrote earlier: I said this had to land with or before the bump. That was wrong — the suite can't be green anyway until mlx releases the operator< fix, so the two are independent.

@PowerBeef

Copy link
Copy Markdown

Sharing notes from a working bump of the vendored core to v0.32.1 on a fork, in case they save this PR a cycle when it moves forward (or forward to 0.32.1):

  • JIT abort at 0.32.1 (does not affect 0.32.0): core 0.32.1's qmv() passes a 6th template argument (results_per_simdgroup, from Optimize large NVFP4 QMV on M5 Max mlx#3961) while the [[kernel]] qmv templates still declare 5 params — with nojit_kernels.cpp excluded (this repo's current configuration), the first quantized matvec aborts at runtime with a Metal front-end "too many template arguments" error. Reported upstream at [BUG] Metal JIT build: qmv template instantiation passes 6 args to 5-param kernel templates (runtime abort) mlx#4350. We switched Cmlx to the nojit path (exclude jit_kernels.cpp instead and resolve kernels from the prebuilt metallib), which sidesteps it.
  • jaccl excludes: the minimal fix from Package.swift exclude list breaks when the mlx pin moves past v0.31.2 (jaccl restructured) #446 (- mesh/ring/utils.cpp, + mlx/mlx/distributed/jaccl/lib) is confirmed working at 0.32.x.
  • mlx-c: the vendored mlx-c needs its compile-cache calls updated for the 0.32.x core API (mlx::core::detail::compile_cache() handle); mlx-c's latest tag (v0.6.0) predates it.
  • tools/update-mlx.sh: the metal kernel make-list needs gemv and searchsorted added at 0.32.x or the regenerated sources fail to link.
  • NAX kernels: with a deployment target below 26.2, the generated *_nax.cpp sources reference symbols the header gates out — we compiled the base build with a MLX_METAL_NO_NAX define and excluded the seven mlx-generated/*_nax.cpp files (a 26.2+ variant can re-enable them).
  • FFT: the C API grew a norm parameter; the Swift FFT.swift call sites need MLX_FFT_NORM_BACKWARD threaded through (~20 sites; the fftshift/ifftshift helpers are unaffected).

Everything above is public at PowerBeef/mlx-swift branch imarello/core-0.32.1 (f2459b6) + PowerBeef/mlx-c branch imarello/core-0.32.1-compat — feel free to lift anything useful.

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).
@davidkoski

Copy link
Copy Markdown
Member

cmake builds will fail until mlx-c gets a tag

@davidkoski

davidkoski commented Aug 29, 2026 •

Copy link
Copy Markdown
Member

OK, not making it today. Next step, debug CUDA build:

terminate called after throwing an instance of 'std::runtime_error'
  what():  cudaStreamSynchronize(stream_) failed: driver shutting down

claude says:

This error means a CUDA API call (cudaStreamSynchronize) was attempted after the CUDA runtime (libcudart) or driver had already begun unloading/shutting down for the process.

It is typically an exit-time teardown issue rather than a failure of the GPU kernel itself.

This is running example1 and indeed it looks like a shutdown error:

  9 [ra] 0x0000729cae8bb384 <unknown> in libstdc++.so.6.0.33
 10 [ra] 0x00005a46c5ced323 mlx::core::check_cuda_error(char const*, cudaError) + 352 in example1
 11 [ra] 0x00005a46c582e09d mlx::core::cu::CommandEncoder::synchronize() + 72 in example1
 12 [ra] 0x00005a46c582c7e8 mlx::core::cu::CommandEncoder::~CommandEncoder() + 27 in example1
 13 [ra] 0x00005a46c583aa3c std::pair<int const, mlx::core::cu::CommandEncoder>::~pair() + 31 in example1
 14 [ra] 0x00005a46c583aa5b void std::destroy_at<std::pair<int const, mlx::core::cu::CommandEncoder> >(std::pair<int const, mlx::core::cu::CommandEncoder>*) + 27 in example1

@davidkoski

Copy link
Copy Markdown
Member

Status: waiting for feedback from mlx folks. It may be a lifecycle issue:

let deviceCtx = mlx_device_new_type(MLX_GPU, 0)
let streamCtx = mlx_stream_new_thread_unsafe(deviceCtx)
mlx_stream_free(streamCtx)
mlx_device_free(deviceCtx)

is a minimal repro. The program crashes when the static std::unordered_map<int, CommandEncoder> encoders gets destroyed on exit.

@davidkoski

davidkoski commented Sep 1, 2026 •

Copy link
Copy Markdown
Member

The CUDA run issue is a new lifecycle requirement: clear_streams must be called on exit. Computation works fine but it crashes on process exit. I am disabling the CUDA runs for now so we can get this merged. There are a couple of PRs that make use of the new mlx. This will not be tagged as a release until mlx-c exposes that API and we adopt it here.

@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. See note on CUDA runs -- waiting for mlx-c updates before tagging release.

@davidkoski
davidkoski merged commit ab924c8 into ml-explore:main Sep 1, 2026
7 checks passed
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>
@davidkoski davidkoski mentioned this pull request Oct 23, 2026
2 tasks done
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.

3 participants