leak global CommandEncoder -- avoid cuda synchronize on process shutdown by davidkoski · Pull Request #4480 · ml-explore/mlx

davidkoski · 2026-09-08T17:57:17Z

used in mlx-swift
global CommandEncoder are still alive after exit even if no work is happening
the destructor tries to call cuda api, which causes a fatal error

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

from

frame #11: 0x0000560c45de753f example1`mlx::core::cu::CommandEncoder::synchronize() + 73
frame #12: 0x0000560c45de5c8a example1`mlx::core::cu::CommandEncoder::~CommandEncoder() + 28
frame #13: 0x0000560c45df3ee4 example1`std::pair<int const, mlx::core::cu::CommandEncoder>::~pair() + 32

☑️ I understand it is strictly prohibited to use AI to write PR description
AI usage disclosure: none used

- used in mlx-swift - global CommandEncoder are still alive after exit even if no work is happening - the destructor tries to call cuda api, which causes a fatal error

davidkoski · 2026-09-08T17:57:59Z

+  // encoders are leaked intentionally as they would synchronize on process
+  // shutdown
+  static auto* encoders = new std::unordered_map<int, CommandEncoder>();
+  return *encoders;


Per internal discussion we decided leaking these on process exit was ok. I was able to test this on a cuda device with mlx-swift.

davidkoski · 2026-09-08T17:58:39Z

  static thread_local std::unordered_map<int, CommandEncoder> encoders;
  return encoders;


I don't know if this has the same issue -- mlx-swift doesn't use the TLS Streams, but since the python code does, perhaps it is ok.

* 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

`~CudaHandle` calls `reset()`, which checks the result through `CHECK_CUDA_ERROR` and throws on failure. A throw that escapes a destructor calls `std::terminate`, so this aborts the process rather than reporting anything. It is reachable at ordinary shutdown. The existing guard tests `cudaPeekAtLastError()`, which catches a sticky per-context error but not the CUDA runtime unloading: once teardown has begun the peek still returns `cudaSuccess`, `reset()` proceeds, `Destroy` returns `cudaErrorCudartUnloading`, and the throw fires. The symptom is a SIGABRT with `Destroy(handle_) failed: driver shutting down` after the program has otherwise finished. The destructor now releases the handle without checking. `reset()` is unchanged, so `operator=` and explicit callers keep their error checking, where throwing is legal. Nothing is lost by not reporting in the destructor: the handle cannot be reclaimed once the runtime is gone. Related: ml-explore#4480 leaked the global `CommandEncoder` map for the same class of teardown problem. The `thread_local` map in `get_command_encoders()` (`device.cpp`) is not leaked, and `Worker` detaches its thread while `~CommandEncoder` only signals `stop()` without waiting, so `CudaStream`, `CudaGraph` and `CudaGraphExec` instances can still be destroyed after the primary context is released. This change makes that safe rather than fatal; whether those owners should also be leaked or joined is a separate question.

* 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>

…tion 3 build-linux-cuda-cmake.sh still checked that cmake/mlx.patch was applied to the fetched MLX. That patch is gone, since both of its fixes ship in v0.32.3 (ml-explore/mlx#4480, #4514), so the check failed every Linux CUDA CMake run right after configure. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

* Fix Log and Equal is_equivalent ignoring primitive state (ml-explore#4266) Co-authored-by: Cheng <git@zcbenz.com> * Stabilize reduced-precision InstanceNorm (ml-explore#4230) * chore: Normalize negative axes in sort and argsort (ml-explore#4332) * Clean up main thread compile cache before python interpreter shuts down (ml-explore#4373) * chore: Check malformed jaccl hostfile that miss rdma in pairs (ml-explore#4284) Co-authored-by: Cheng <git@zcbenz.com> * Round mxfp8 block scales up to avoid saturation (ml-explore#4353) Co-authored-by: Daniel Hiltgen <daniel.hiltgen@ollama.com> Co-authored-by: Cheng <git@zcbenz.com> * Add support for the __array_namespace_info__ (ml-explore#4334) * Stop a failed CUDA graph commit from poisoning the encoder (ml-explore#4356) Co-authored-by: Cheng <git@zcbenz.com> * Fix quantized kernels in JIT build (ml-explore#4372) Co-authored-by: Cheng <git@zcbenz.com> * Avoid zero work in stride-2 ConvTranspose3d (ml-explore#4343) * [CUDA] Ce fused kernel (ml-explore#3947) * Fix cpu exclusive scan for complex numbers (ml-explore#4272) Co-authored-by: Cheng <git@zcbenz.com> * Support Relocatable CUDA DLLs on Windows (ml-explore#4382) * Use cast_to for fused AsType in compiled Metal kernels (ml-explore#4351) Co-authored-by: katlun-lgtm <katlun@windyviews.com> Co-authored-by: Cheng <zcbenz@gmail.com> * python: Declare DLPackCompatible protocol members as methods (ml-explore#4384) * Fix quantizing sliced arrays (ml-explore#4381) * Fix einsum dropping a trailing empty subscript (ml-explore#4299) Co-authored-by: Cheng <git@zcbenz.com> * Add script to run python tests (ml-explore#4393) * Hold GIL in AttachedData destructor (ml-explore#4391) * Patch bump to 0.32.3 (ml-explore#4406) * Fix scan and sort for a zero-size axis (ml-explore#4340) * python: Add a correction parameter in std and var (ml-explore#4348) Signed-off-by: Pradyot Ranjan <99216956+pradyotRanjan@users.noreply.github.com> Co-authored-by: Pradyot Ranjan <99216956+pradyotRanjan@users.noreply.github.com> Co-authored-by: Cheng <git@zcbenz.com> * Fix tests/run.py stuck in macOS CI (ml-explore#4410) * chore: Update extension example (ml-explore#4412) * Fix sorted gather_qmm NAX row overflow above 32K (ml-explore#3922) Co-authored-by: Cheng <git@zcbenz.com> * Fix deadlock caused by mx.clear_streams() holding GIL (ml-explore#4413) * Fix integer pow zeroing a whole SIMD vector on a negative exponent (ml-explore#4354) * Make concurrency cap on load adaptative so that I/O scales with the machine (ml-explore#4408) * Fix pad with an axes subset and negative axes (ml-explore#4364) * Use sdpa_vector_2pass_1_gqa for GQA size 12 and 16 (ml-explore#4380) * python: Make axis of put_along_axis optional (ml-explore#4360) * python: Make axis default to -1 in take_along_axis (ml-explore#4368) * Propagate NaN in arg reductions (ml-explore#4291) Co-authored-by: Cheng <git@zcbenz.com> * Fix mx.from_fp8 E4M3FN NaN decode with branchless carry (ml-explore#4376) Co-authored-by: Saud <saudkamran@Sauds-MacBook-Pro.local> * Update rules for PR limitation bypass list (ml-explore#4415) * [CUDA] Use native events for GPU fence waits (ml-explore#4401) Co-authored-by: Cheng <git@zcbenz.com> * Validate GGUF tensor dimensions (ml-explore#4378) Co-authored-by: Cheng <git@zcbenz.com> * Fix crash on ellipsis indexing with too many trailing indices (ml-explore#4396) Co-authored-by: Anastasiia Filippova <a_filippova@apple.com> * [CUDA] Cholesky via cuSOLVER (ml-explore#4208) Co-authored-by: Cheng <git@zcbenz.com> * python: Fix __setitem__ with negative index preceded by None (ml-explore#4414) * Update nanobind to 3.0.1 (ml-explore#4417) * chore: Use normalize_axis_index in scan ops (ml-explore#4383) * python: Add explicit index and bytes support (ml-explore#4388) * docs: fix typo recived -> received (ml-explore#4424) Signed-off-by: Vaibhav Srivastava <vaibhavsri1712@gmail.com> * Fix Metal grid sizing for strided scans (ml-explore#4420) Co-authored-by: Cheng <git@zcbenz.com> * Fix SGD and Adafactor weight decay mutating caller arrays (ml-explore#4426) * Widen float16/bfloat16 to float32 in cpu reduce (ml-explore#4387) * Fix sdpa_vector_2pass_1_gqa kernel batch offset for K/V in gqa decode (ml-explore#4431) * Make head-dim-256 prefill with array mask run fused NAX kernel (ml-explore#4416) * python: Make unstack return tuple instead of list (ml-explore#4448) * python: Make iter() throw TypeError for 0-dim array (ml-explore#4425) Co-authored-by: Cheng <git@zcbenz.com> * Fix save and save_safetensors corrupting a lazily loaded source file (ml-explore#4434) * [Metal] Use hypot for complex ops (ml-explore#4345) Co-authored-by: Cheng <git@zcbenz.com> * Add get_array_buffer_size to query buffer size for arrays (ml-explore#4436) * Add M5 ultra tunings for non-quantized matmuls (ml-explore#4447) * [CUDA] Fix completion worker busy loop (ml-explore#4452) * Fix non transposed affine qmm dispatch logic (ml-explore#4392) * Fix sorted gather_qmm on ragged K (ml-explore#4009) Co-authored-by: Cheng <git@zcbenz.com> * Don't restore a thread-affine stream when StreamContext dies on another thread (ml-explore#4462) * Fix Pad vjp with an axes subset and negative axes (ml-explore#4441) Co-authored-by: Cheng <git@zcbenz.com> * Fix integer div-by-zero crash in CPU backend (sort, scan, binary ops) (ml-explore#4442) Co-authored-by: Cheng <git@zcbenz.com> * Use 64-bit file seeks on Windows (ml-explore#4456) Co-authored-by: Cheng <git@zcbenz.com> * Use precise::exp in Sigmoid so compiled and eager sigmoid agree (ml-explore#4461) Co-authored-by: Pierre Lamy <pierrelamy427@gmail.com> * Use NAX attention for long unmasked D72/D80 inputs (ml-explore#4455) * Add D512 support to Metal vector attention (ml-explore#4459) Co-authored-by: Cheng <git@zcbenz.com> * Add a scoped_env utility for tests (ml-explore#4474) * [Metal] global scale for qmm (ml-explore#4458) * Fix re-entering the same stream context manager (ml-explore#4478) * Leak global CommandEncoder on exit (ml-explore#4480) * Update CONTRIBUTING.md (ml-explore#4469) * python: Fix bytes() on non-contiguous arrays (ml-explore#4449) * Break the sibling cycle when an array is released by assignment (ml-explore#4453) Co-authored-by: Cheng <git@zcbenz.com> * win: Use DXGI for accurate WDDM VRAM memory budget (ml-explore#4457) Co-authored-by: Cheng <git@zcbenz.com> * Add type hints and docstrings for ALiBi layer (ml-explore#4472) * Fix fp quantized matmul corruption when the quantized dim is not a multiple of 32 (ml-explore#3912) Co-authored-by: Rohan Gautam <rohan1gautam@gmail.com> * Fix CUDA test synchronization flake (ml-explore#4490) * ci: Skip bypass list updates in forks (ml-explore#4495) * Load global scales in qmm_t kernels (ml-explore#4483) * python: Release the GIL when copying Metal DLPack inputs (ml-explore#4497) Signed-off-by: Ranran Haoran Zhang <ranranhaoranzhang@gmail.com> * Use matrix kernels for global-scale gather_qqmm (ml-explore#4481) * Adding metal kernels for the gated delta nets. (ml-explore#4020) Co-authored-by: Yanzhao Wang <19340816+wyanzhao@users.noreply.github.com> Co-authored-by: Cheng <git@zcbenz.com> * Fix jaccl ring all_gather: direction 1 slice was not mirrored (ml-explore#4443) Co-authored-by: Dave Hind <davehind42@gmail.com> Co-authored-by: Anastasiia Filippova <a_filippova@apple.com> * Skip Metal-only gated delta kernel tests on other backends(fix CI) (ml-explore#4522) Co-authored-by: Anastasiia Filippova <a_filippova@apple.com> * [CUDA] Add global scale support to gather_qmm (ml-explore#4507) Co-authored-by: Anastasiia Filippova <a_filippova@apple.com> * [BUG][Metal] Deadlock in fence (ml-explore#4552) Co-authored-by: Cheng <zcbenz@gmail.com> * Fix ordering of complex64 with NaNs (ml-explore#4519) * Improve metal memory usage for SDPA D256 (ml-explore#4505) Co-authored-by: Cheng <git@zcbenz.com> * Add support for matrix transpose (ml-explore#4402) Co-authored-by: Cheng <git@zcbenz.com> * Improve metal memory usage for SDPA D512 (ml-explore#4487) * Support shapeless compilation of scan operations (ml-explore#4510) * [Metal] Gather mm improvement (ml-explore#4567) * Fix perf regression for upsample `linear` mode align_corners=False (ml-explore#4500) * Add fused Metal kernels for fast.cross_entropy (ml-explore#4520) Co-authored-by: Anastasiia Filippova <a_filippova@apple.com> * python: Add missing defaults to tri, tril, triu, gather_mm signatures (ml-explore#4492) * Add Metal SDPA support for D96/V64 (ml-explore#4499) Co-authored-by: Cheng <git@zcbenz.com> * Raise IndexError for out of bounds axes (ml-explore#4484) * Fix floor_divide for integers (ml-explore#4515) Co-authored-by: Cheng <git@zcbenz.com> * Fix cpu binary ops on large data (ml-explore#4517) Co-authored-by: Cheng <git@zcbenz.com> * Do not throw from CUDA destructors and avoid implicit default streams in eval/compile (ml-explore#4514) Co-authored-by: Cheng <git@zcbenz.com> * [Metal] gather_qmm improvement (ml-explore#4572) * Fix floating-point constant precision in compiled kernels (ml-explore#4511) * chore: Fix assertion in ReduceScatter::eval_cpu (ml-explore#4557) * Leak thread local streams on exit in main thread (ml-explore#4576) * Adding VJP for the scaled dot product attention. (ml-explore#4563) * python: Add __pos__ support to array (ml-explore#4488) Co-authored-by: Cheng <zcbenz@gmail.com> * Add Python 3.15 and 3.15t support (ml-explore#4555) Co-authored-by: Cheng <git@zcbenz.com> * Add fixed size `unique` operation (ml-explore#4501) Co-authored-by: Cheng <zcbenz@gmail.com> * Fix normalizing single-element slice with negative stride (ml-explore#4446) Co-authored-by: Cheng <git@zcbenz.com> * Patch bump to 0.32.4 (ml-explore#4582) * Adding VJP for gated delta net. (ml-explore#4565) * Fix `col_reduce_longcolumn` for negative-stride views (ml-explore#4529) Co-authored-by: Cheng <zcbenz@gmail.com> * Clear CUDA errors once reported and raise eval errors in buffer protocol (ml-explore#4523) Co-authored-by: Cheng <zcbenz@gmail.com> * Improve pre-M5 metal memory usage for SDPA D512 (ml-explore#4518) * Fix split data size span (ml-explore#4540) * Re-introduce Array API Tests suite (ml-explore#4433) Signed-off-by: Pradyot Ranjan <99216956+pradyotRanjan@users.noreply.github.com> Co-authored-by: Pradyot Ranjan <99216956+pradyotRanjan@users.noreply.github.com> Co-authored-by: Cheng <zcbenz@gmail.com> * Fix view inside of a vmap (ml-explore#4586) * Fix cumsum of a non-contiguous bool (ml-explore#4590) * docs: contraction dtype requirements in matmul/einsum (ml-explore#4550) Co-authored-by: Cheng <zcbenz@gmail.com> * Fixing cooperative tensor build issues on Mac OS 27 (ml-explore#4594) * Handle EOF and EINTR in ParallelFileReader::read (ml-explore#4532) Co-authored-by: Cheng <git@zcbenz.com> * Clamp slice bounds when a slice is combined with an array index (ml-explore#4513) Co-authored-by: Cheng <zcbenz@gmail.com> * Fix grouped Conv1d/Conv2d init to use in_channels // groups (ml-explore#4544) Co-authored-by: Cheng <zcbenz@gmail.com> * Fix mx.view misaligned offset (ml-explore#4585) Co-authored-by: Cheng <git@zcbenz.com> * [CUDA] Add fp_gather_qmv to optimize gather_qmm (ml-explore#4508) * Fix eval cleanup deadlock by committing GPU streams before waiting (ml-explore#4556) * Fix shared_buffer_reshape contiguity flags (ml-explore#4589) * Fix view last axis stride (ml-explore#4587) Co-authored-by: Cheng <zcbenz@gmail.com> * Fix empty unsigned sum/prod failing on Metal (ml-explore#4546) Co-authored-by: Cheng <git@zcbenz.com> * [CUDA] Add gather_qmm_rhs_sm80 (ml-explore#4554) Co-authored-by: Cheng <git@zcbenz.com> * python: Fix bare ellipsis __setitem__ by squeezing leading singleton dimensions (ml-explore#4561) * Fix row/col contiguous flags in view (ml-explore#4588) * Do not cache failed distributed backend initialization (ml-explore#4558) * Fix reduction over large arrays (ml-explore#4549) Co-authored-by: Cheng <git@zcbenz.com> * docs: Correct the export callback type in the guide (ml-explore#4578) * Fix MultiOptimizer crash when an optimizer's group is empty (ml-explore#4570) * Fix fft size one (ml-explore#4591) * Sanitize '-' in make_template_hash for metal kernel (ml-explore#4584) Signed-off-by: GokayAI <60583610+gokay-ai@users.noreply.github.com> Co-authored-by: Cheng <git@zcbenz.com> * Fix as_strided contiguity (ml-explore#4575) * Reject non-aligned pointers when converting to mx.array (ml-explore#4592) * Fix conjugation in complex VJPs for inverse trig ops, cos, sqrt, and rsqrt (ml-explore#4606) * Fix complex to bool (ml-explore#4593) * [jaccl] second cable for the mesh (ml-explore#4607) * Fix vmap of scatter placing the batch axis wrongly in the updates (ml-explore#4322) Co-authored-by: Cheng <git@zcbenz.com> * Unroll TK=4 keys per iter in sdpa_vector_2pass_1 (ml-explore#4596) Co-authored-by: Cheng <git@zcbenz.com> * Fix CPU float16 SIMD comparison masks (ml-explore#4602) Co-authored-by: Cheng <git@zcbenz.com> * Add jvp support for logcumsumexp (ml-explore#4611) * [CUDA] Export CommandEncoder and Device (ml-explore#4619) * Use the fast qmv kernel for outputs not divisible by 8 (ml-explore#4516) Co-authored-by: Cheng <git@zcbenz.com> * Fix ResidencySets reading autoreleased error object (ml-explore#4623) * Use the small-channel implicit gemm loader for C_per_group = 8 (ml-explore#4617) * [CUDA] Reduce floor for fp_gather_qmv (ml-explore#4620) * Fix default.metallib lookup under swift test with Swift Build (ml-explore#4527) * Fix metal validation error when passing empty vector (ml-explore#4622) Co-authored-by: Cheng <git@zcbenz.com> * Replace Black formatter with ruff (ml-explore#4600) * Tune M1 quantized matmul dispatch for six-row large shapes (ml-explore#4629) Co-authored-by: deadbydawn <gothravenllm@gmail.com> Co-authored-by: Cheng <git@zcbenz.com> * [metal] Add env var based residency refresh interval (ml-explore#4633) * Promote integer inputs in logcumsumexp (ml-explore#4625) Co-authored-by: Cheng <git@zcbenz.com> * Fix slicing's non contiguous start (ml-explore#4599) * Fix float16 overflow in clip_grad_norm (ml-explore#4626) * Collapse the non sorted axes to pick the contiguous sort kernel (ml-explore#4366) Co-authored-by: Cheng <git@zcbenz.com> * Fix 32 bit dispatch grid overflow in Metal reductions (ml-explore#4437) Co-authored-by: Cheng <git@zcbenz.com> * Fix gpu boolean masking in 0 dim array (ml-explore#4635) Co-authored-by: Cheng <git@zcbenz.com> * Fix vmap of an inverse real FFT to an odd length (ml-explore#4637) * Fix cross-process race in CPU compile cache (ml-explore#4638) Co-authored-by: Cheng <git@zcbenz.com> * Update acknowledgments for past pinv contribution (ml-explore#4645) * Fix foreign buffer ownership (ml-explore#4634) * Add MLX_CONV_WINOGRAD switch to disable winograd for conv2d (ml-explore#4639) * Keep split-K quantized matmul partials in float32 (ml-explore#4641) * Thin N GEMM kernels (ml-explore#4640) Co-authored-by: Cheng <git@zcbenz.com> * Add 1-bit affine quantization support (Metal) (ml-explore#3161) Co-authored-by: Angelos Katharopoulos <a_katharopoulos@apple.com> Co-authored-by: bri-prism <288398250+bri-prism@users.noreply.github.com> * Update Metal Complex division in numpy style to avoid overflow condition (ml-explore#4644) * Fix pickle bfloat16 strides (ml-explore#4649) * Fix metal negative stride index in masked_scatter/scatter/SDPA mask/segmented_gemm (ml-explore#4650) * Fix missing break binary bool (ml-explore#4651) * Fix cpu binary_op_dispatch_dims on large data (ml-explore#4648) * Add FloorDivide primitive for integer floor division (ml-explore#4642) Co-authored-by: Cheng <git@zcbenz.com> * [Metal] Fix SDPA max threadgroup for M1/M2 generation (ml-explore#4643) Co-authored-by: Cheng <git@zcbenz.com> * Fix incorrect output mask computation in `mx.block_masked_mm` (ml-explore#4652) * Add thin M GEMM dispatch (ml-explore#4654) * Fix Python dtpye -> mlx conversion narrowing (ml-explore#4656) * Fix input/output scalar merging in compile (ml-explore#4658) Co-authored-by: Cheng <git@zcbenz.com> * Validate .npy shape in mlx/io/load.cpp (ml-explore#4657) Co-authored-by: Cheng <git@zcbenz.com> * Replace "pull_request_target" event with "pull_request" in update_bypass_list workflow (ml-explore#4667) * [CUDA] Only count consecutive misses in the LRUCache thrashing check (ml-explore#4661) Co-authored-by: Cheng <git@zcbenz.com> * Fix conv input gradient shape with uneven or large padding (ml-explore#4662) * Fix missing stream in layer_norm and sdpa fallbacks (ml-explore#4663) Co-authored-by: Cheng <git@zcbenz.com> * Fix direct CPU conv writing past output (ml-explore#4668) Co-authored-by: Cheng <git@zcbenz.com> * Add stft and istft to mlx.core.fft (ml-explore#3639) * Fix lost rank output in the distributed launcher (ml-explore#4615) Co-authored-by: Cheng <git@zcbenz.com> * Fix split-K test compatibility on M5, M6 (ml-explore#4677) * Fix signedness warning in quantized Metal tail bounds check (ml-explore#4675) * Fix batched row count mismatch in metal qvm_split_k (ml-explore#4673) Co-authored-by: Cheng <git@zcbenz.com> * Fix metal conv bounds checks and integer overflows (ml-explore#4672) Co-authored-by: Cheng <git@zcbenz.com> * python: Make arctan2/radians/degrees support scalar inputs (ml-explore#4679) Co-authored-by: Cheng <git@zcbenz.com> * Fix missing add_temporary in pad_gpu/fill_gpu (ml-explore#4660) Co-authored-by: “Ronald <“Ronaldmannak@me.com”> * Enable non-transposed NAX gather_qmm_rhs (ml-explore#4684) * Add Metal extension build support (ml-explore#4004) * Adding minmax method to MLX to follow Numpy/Jax API (ml-explore#4674) Co-authored-by: Cheng <git@zcbenz.com> * Fix missing stream in gather matmul indices cast (ml-explore#4687) * [BUG] Metal fence deadlock (ml-explore#4689) * Patch bump to 0.32.5 (ml-explore#4707) * docs: Publishing a New Release (ml-explore#4691) * ci: Run update_bypass_list on push instead of pull_request (ml-explore#4671) * [CUDA] Run the vector SDPA kernels at head_dim 256 (ml-explore#4698) Co-authored-by: yanghaojin <haojin.yang@greenbit.ai> * [CUDA] Fix grouped gemm with large data (ml-explore#4696) Co-authored-by: yanghaojin <haojin.yang@greenbit.ai> Co-authored-by: Cheng <git@zcbenz.com> * Fix conv2d OOB reads for partial output-channel tiles (ml-explore#4653) * docs: MetalLibrary in custom extension guide (ml-explore#4701) * [Metal] Fix SDPA f16/f32 d512 max threadgroup for M1/M2 generation (ml-explore#4708) * Add float16 head_dim 256 TK=1 SDPA vector kernel for M1/M2 (ml-explore#4719) Co-authored-by: Cheng <git@zcbenz.com> * Allow configuring the CPU JIT cache directory (ml-explore#4714) Co-authored-by: Cheng <git@zcbenz.com> * Make `broadcast_arrays` return a tuple (ml-explore#4713) Co-authored-by: Cheng <git@zcbenz.com> * Fix unaligned scalar reads and oversized contiguous buffer views (ml-explore#4721) Co-authored-by: Cheng <git@zcbenz.com> * [Metal] Add staged affine_qmm_t_nax kernels with split-k (ml-explore#4682) Signed-off-by: Yanzhao Wang <yzwang1996@gmail.com> Co-authored-by: Cheng <git@zcbenz.com> * [Metal] Add fast route for dequantizing fp4 to half in fp_qmm_t (ml-explore#4709) Co-authored-by: Cheng <git@zcbenz.com> * [Metal] Add qmv_mma/qmv_nax kernels for 6-16 rows with large N/K (ml-explore#4712) Co-authored-by: Cheng <git@zcbenz.com> * docs(forkdiff): move the base to upstream c6abf0d after the merge Set base.hash to the new merge-base. Remove the v0.32.2 sync section and the gather_mm row-tile backport section, because upstream now has those changes. Keep the revert of the ml-explore#4352 simdgroup skip in its own section. Reduce the floor_divide section to the tests the fork keeps. Add the re-deleted bypass-list workflow and scripts to Fork tooling. Keep the revert of the ml-explore#4352 simdgroup skip in fp_gather_qmm_rhs_nax and affine_gather_qmm_rhs_nax. The merge had taken upstream's sg_active guard there. State the new upstream sync policy in FORKDIFF.md: merge upstream when David directs it, and move base.hash at each merge. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * fix: repair the merge in the quantized kernels and the lint The merge put upstream's new "typename U = T" template parameter on the fork's qmm_t_expert_impl instead of qmm_t_impl. Move it back. fp_gather_qmv_fast_tail now calls adjust_matrix_offsets with the global scale, as upstream's fp_gather_qmv_fast does. Remove the extra closing brace that the previous commit put in affine_gather_qmm_rhs_nax. The kernel body is again the fork's body. Import numpy in test_fast_sdpa.py again. Upstream removed the import, and the fork's FP32-partials tests use it. Wrap the long lines in check_forkdiff.py for the ruff E501 rule that upstream turned on. Remove the PR #15 bias-sum sentence from fork.yaml: upstream now promotes the bias-sum operands itself (ml-explore#4516). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * test: follow the upstream allocator and gather_qmm APIs in fork tests Upstream allocator::malloc now returns an owning allocator::Data, and allocator::free is gone. The fork's snapshot and footprint tests need explicit release points, so they use allocator::allocator().malloc and free, as upstream's own allocator tests do. Upstream gather_qmm takes a global_scale argument before sorted_indices. Pass std::nullopt in the GPT-OSS MXFP4 tests. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * fix(cpu): floor_divide of min by -1 must not trap Upstream's FloorDivide divides the minimum signed value by -1 on the CPU. The quotient overflows, and on x86 the division traps (SIGFPE). The fork test "test floor_divide integer edge cases on the CPU" crashed in the Fedora x86_64 job. Divide by 1 there instead, as the PR #32 functors did. The result wraps to min. Describe the guard in the floor_divide section of fork.yaml. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * Compute the JACCL ring wire offsets in 64 bits CodeQL (cpp/integer-multiplication-cast-to-long) flagged the ring's wire offset expressions on PR 38: `lr * n_wires` is an `int` product that is widened to 64 bits only after the multiplication. The same pattern is in three ring_impl.h loops and in the two `size_ * 2 * n_conns_` thresholds in ring.cpp. Widen the first operand to int64_t so every product is computed in 64 bits. The operands are small, so results do not change. fork.yaml describes the new fork change. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> * Compute the vector SDPA query layout check in 64 bits CodeQL (cpp/integer-multiplication-cast-to-long, alert 29) fails the PR 38 check on the vector SDPA query layout test: `shape[3] * shape[bidx]` is an `int` product compared with a 64-bit stride. Widen `shape[3]` to int64_t. The shapes are small, so the result does not change. The fork.yaml section for the JACCL ring fix now covers this file too. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> * Fix `mx.take`'s negative indices (ml-explore#4729) * Fix normal loc scale shape (ml-explore#4730) * Fix zero-copy numpy arrays dtype mismatch (ml-explore#4731) * ci: skip the 4 GiB reduce test where the runner cannot hold it Upstream's test_init_reduce_large_output (ml-explore#4437) makes a 4 GiB boolean output. The fork CI runs on a GitHub-hosted macOS runner with a 3.5 GiB Metal buffer limit and 7 GB of memory. On the GPU the allocation fails; on the CPU it competes with the other test files for memory. Skip the test when the default device cannot hold the output. Machines with more memory still run it. When a CPU test file fails, print the end of its log at once. The job of the previous run was stopped before the full logs printed, so the cause of the test_eval failure on the CPU is not known. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * Widen affine qmv_wide row tiles to 8 (ml-explore#4728) * Fix/padding jvp vjp (ml-explore#4736) * fix: qqmm refuses a global scale on the CPU; adapt GPU-path tests The CPU QQMatmul kernel has no global scale input and ignores the scales that qqmm passes to it. qqmm with global scales on a CPU stream returned a wrong result with no error: in the fork CPU test pass, test_fp_qmv_fast_unaligned_n_global_scale differed by about 1000. qqmm now raises "Global scale is only supported on the GPU" on a stream that is not a GPU stream, as gather_qmm already does. Upstream tests that check only GPU kernels failed in the CPU pass: test_sdpa_vector_2pass_threadgroup_limit (force_fused=True), test_fp_qmv_fast_unaligned_n_global_scale, test_qqmm_global_scale_matrix_kernels and test_gather_qmm_global_scale. When the default device is the CPU, they now check that the documented error is raised. The GPU pass runs them as before. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * ci: add a temporary diagnostic for the CPU NaN; reword the 4 GiB skip Add one temporary step after the build in metal-build-and-test. It runs the failing case of test_fp_gather_qmm_non_multiple_of_32 (K=512, N=1040, transpose=False, sorted, float32) on the CPU with a fresh allocator, after a warm-up that fills the allocator cache with NaN buffers, and with the buffer cache off, and on the GPU. It prints NaN counts per expert for w, w_hat, y_q and y_hat. The step has continue-on-error, so the full test suite still runs. The next push removes it. State in fork.yaml when test_init_reduce_large_output skips, why it cannot be smaller (ml-explore#4437), and that David ruled on 2026-10-23 to keep the skip. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * ci: move the CPU NaN diagnostic to its own short job Remove the temporary diagnostic step from metal-build-and-test. Its reimplementation found no NaN in any mode. Add a temporary job, diag-cpu-nan, that builds the wheel and runs only the diagnostic, so its log is ready after the build. It runs the real test_fp_gather_qmm_non_multiple_of_32 alone, then after the tests that come before it, and an instrumented copy of the test body (only the failing case, and all cases in the test order, each evaluated as one graph). The findings also go to ::notice annotations. The next push removes the job. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * ci: install numpy for the temporary NaN diagnostic job The tests import numpy through mlx_tests. The diagnostic job installed only the wheel, so every run stopped at the import. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * Fix data race in distributed ring on pending send/recv lists (ml-explore#4737) Co-authored-by: Cheng <git@zcbenz.com> * Use data_size for gather/scatter large-offset checks (ml-explore#4724) Co-authored-by: Cheng <git@zcbenz.com> * fix(cpu): clear C before the BLAS matmul when beta is 0 test_fp_gather_qmm_non_multiple_of_32 (K=512, N=1040, sorted, float32) gave NaN on the CPU. The diagnostic job found the NaN only in the reference y_hat = x @ w_sel, a batched matmul with M = 1. Its inputs had no NaN. The test failed alone too, but only after the earlier subtests of the same test: y_hat is 266240 bytes, the same size as the packed fp4 weights of the previous subtest, so the output reused that buffer from the allocator cache. Accelerate read C even with beta = 0 and computed NaN * 0. Fill C with zeros before cblas_sgemm, cblas_dgemm and cblas_cgemm when beta is 0. MLX's own SIMD gemm already skips C when beta is 0. Remove the temporary diagnostic job. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * fix(cpu): clear C at every BLAS gemm call with beta 0 The previous commit cleared C only in the gemms/cblas.cpp wrappers. Four more call sites call cblas_sgemm directly with beta 0: BlockMaskedMM and GatherMM (the CPU MoE gathered matmul) in masked_mm.cpp, and the two explicit-gemm convolutions in conv.cpp. They can give the same NaN * 0 from a reused output buffer. Add one helper, clear_gemm_output in gemm.h, that fills the exact C region of a call (rows by columns, with its ldc) with zeros. Use it at all five call sites, the wrappers included. No other BLAS gemm or gemv call exists in the CPU backend. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * ci: build once and run the Metal and CPU test passes in parallel Split the fork-tests job into parallel jobs. build-wheel and build-cpp build in parallel. build-cpp also runs the C++ tests on the GPU and on the CPU. python-gpu (3 shards) and python-cpu (4 shards) run the Python tests from the uploaded wheel. metal-build-and-test merges the coverage profiles of all jobs, writes the same coverage report, and fails when another job failed. python_shards.py checks the shard map: each python/tests/test*.py file must be in each pass exactly one time. Split files must partition their test IDs. The last job fails when a test ran in two shards. fork-build-cache restores and saves ccache on pull requests and on main, and keeps a cache of the compiled Metal kernels through an xcrun wrapper keyed on the preprocessed source, the arguments and the toolchain. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> * [CUDA] Fix gather_qmm for experts past 2^32 weight elements (ml-explore#4738) Co-authored-by: Cheng <git@zcbenz.com> * ci: split the long CPU test files and skip cache saves with no misses Split test_fast_sdpa and test_quantized into named tests and rest entries in the CPU shards, so that no CPU shard has an entry longer than test_gated_delta_sequential. Put the entries in longest-first order. Save the ccache files and the Metal kernel cache only when the build had misses. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> * ci: key the build caches on compiler versions, not the runner image The full --version text of clang and metal names the install folder and the OS, and it changed between three macos-26 runner images, so every job started from an empty cache. Key ccache on the clang major.minor version and fall back to a cache of another clang version; ccache keeps compiler_check=content. Key the Metal kernel cache on the metal compiler version and the kernels hash, and fall back only to a cache of the same metal compiler. Each Metal entry key holds the first line of the metal --version text and the SDK build. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> * Add array::is_uniquely_referenced for the fork C API Upstream 5dbe999 (ml-explore#4634) makes is_donatable() also require that MLX owns the buffer, so the allocator never reuses a wrapped external buffer. The fork's mlx-c buffer-info API reported `unique` from is_donatable(), and mlx-swift-lm uses that as a "no other reference" proof for paged KV segments, checkpoint memory accounting and MTP state, including wrapped host pages, which now always read as not unique. Add is_uniquely_referenced(): the same reference-count test without the ownership test. is_donatable() and donation behavior do not change. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> * fix: async_eval must not throw a stream task error test_eval.test_async_eval_error_in_synchronize (upstream ml-explore#4338) failed on some runs of the fork's CPU test pass: mx.async_eval raised the LU factorization error that the test expects only from mx.synchronize. Cause: StreamThread::enqueue throws the stream's stored error when it runs on the main thread (ml-explore#3742). async_eval enqueues several tasks. When the stream thread finishes the failing task before async_eval enqueues the next one, that enqueue throws. The order depends on thread timing, so the failure is intermittent. The fork does not change this path. async_eval now holds a scheduler::DeferStreamErrors guard. While it lives, enqueue on that thread does not throw the stored error, and synchronize throws it later. eval does not change. Add a C++ test that stores the error before async_eval, so it fails on every run without the fix. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * Fix GPU pad for outputs larger than 2^31 elements (ml-explore#4744) * ci: format python_shards.py with the pinned ruff Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> * ci: run the Python tests in 2 GPU and 3 CPU shards The hosted macOS runners are a shared pool. Pack the GPU pass into 2 shards and the CPU pass into 3 shards, balanced by the durations that run 39902948153 measured. A run now uses 8 macOS jobs, and at most 6 run at the same time (before: 10 jobs, 8 at the same time). Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> * ci: give the NAX kernels Metal cache hits and lighten CPU shard 0 The preprocessed source of the NAX kernels names the Metal toolchain mount folder, which has a random suffix on each runner. The 10 NAX kernels missed the Metal cache in every warm build (41 of 51 hits). Remove the suffix from the text that the key hashes, and keep the toolchain version. CPU shard 0 was the slowest shard in the cold and warm runs of b72420d (651 s and 470 s of tests). Move test_quantized:rest to CPU shard 1 (221 s). Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> * docs(forkdiff): move the base to upstream afc03c5 Set base.hash to the new merge-base after the merge of the 9 upstream commits. No fork section changes: the merge has no conflict, and each merged file gets exactly upstream's change. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> --------- Signed-off-by: Pradyot Ranjan <99216956+pradyotRanjan@users.noreply.github.com> Signed-off-by: Vaibhav Srivastava <vaibhavsri1712@gmail.com> Signed-off-by: Ranran Haoran Zhang <ranranhaoranzhang@gmail.com> Signed-off-by: GokayAI <60583610+gokay-ai@users.noreply.github.com> Signed-off-by: Yanzhao Wang <yzwang1996@gmail.com> Co-authored-by: rohith <kapellirohith@gmail.com> Co-authored-by: Cheng <git@zcbenz.com> Co-authored-by: Vladimir Iglovikov <ternaus@users.noreply.github.com> Co-authored-by: Brian C. <94733710+deBrian07@users.noreply.github.com> Co-authored-by: Cheng <zcbenz@gmail.com> Co-authored-by: Erwin Zhang <59893706+erwinzhang7@users.noreply.github.com> Co-authored-by: Daniel Hiltgen <dhiltgen@users.noreply.github.com> Co-authored-by: Daniel Hiltgen <daniel.hiltgen@ollama.com> Co-authored-by: Aaishwarya Mishra <aaishwarymishra@gmail.com> Co-authored-by: YH Yan <strayberry0w0@gmail.com> Co-authored-by: Dwijen Patel <dwijen@gmail.com> Co-authored-by: Anastasiia Filippova <a_filippova@apple.com> Co-authored-by: Ayaan Gazali <ayaangazali.work@gmail.com> Co-authored-by: katlun-lgtm <katlun@gmail.com> Co-authored-by: katlun-lgtm <katlun@windyviews.com> Co-authored-by: Adityaj0 <93090622+Adityaj0@users.noreply.github.com> Co-authored-by: Eyüp Can Akman <eyupcanakman@gmail.com> Co-authored-by: Pradyot Ranjan <99216956+prady0t@users.noreply.github.com> Co-authored-by: Pradyot Ranjan <99216956+pradyotRanjan@users.noreply.github.com> Co-authored-by: Jagrit Digani <digani@apple.com> Co-authored-by: Philip John Basile <PBasile@Basilecom.com> Co-authored-by: Alessio Pollero <alessio.pollero@ventureone.ae> Co-authored-by: Duhyeon, Kim <49020301+dudududukim@users.noreply.github.com> Co-authored-by: Atirna <noctiselara0@gmail.com> Co-authored-by: Saud Kamran <32072820+saud5150@users.noreply.github.com> Co-authored-by: Saud <saudkamran@Sauds-MacBook-Pro.local> Co-authored-by: YH Yan <yh.yukiyan@gmail.com> Co-authored-by: Roshan Sharma <rs127@fordham.edu> Co-authored-by: sashko-zakharchuk <oleksandr.zakharchuk@gmail.com> Co-authored-by: XXXXRT666 <157766680+XXXXRT666@users.noreply.github.com> Co-authored-by: Vaibhav Srivastava <vaibhavsri1712@gmail.com> Co-authored-by: Robin Gieseke <38696388+TheDarkchip@users.noreply.github.com> Co-authored-by: 2sumtech <2sumtech@gmail.com> Co-authored-by: Vedant Singhal <88609765+Ved235@users.noreply.github.com> Co-authored-by: JasonHonKL <148705846+JasonHonKL@users.noreply.github.com> Co-authored-by: Simeet Nayan <78461155+simeetnayan81@users.noreply.github.com> Co-authored-by: Cdimoy <50119145+Cdimoy@users.noreply.github.com> Co-authored-by: Alessio Pollero <alessio.pollero@gmail.com> Co-authored-by: Rohan Gautam <rohan1gautam@gmail.com> Co-authored-by: michalk8 <46717574+michalk8@users.noreply.github.com> Co-authored-by: Nilesh Patil <128893479+nileshpatil6@users.noreply.github.com> Co-authored-by: Marcos Ashton <59975445+MarcosAsh@users.noreply.github.com> Co-authored-by: pierre427 <pierre@userid.org> Co-authored-by: Pierre Lamy <pierrelamy427@gmail.com> Co-authored-by: Yanzhao Wang <19340816+wyanzhao@users.noreply.github.com> Co-authored-by: David Koski <46639364+davidkoski@users.noreply.github.com> Co-authored-by: Isaac Hernández <isadohergar@gmail.com> Co-authored-by: Tudorica Constantin-Alexandru <tudalex@gmail.com> Co-authored-by: Ritaban Mitra <39273308+Ritabanm@users.noreply.github.com> Co-authored-by: Ranran <ranranhaoranzhang@gmail.com> Co-authored-by: Tommaso Pegolotti <31819389+tpegolotti@users.noreply.github.com> Co-authored-by: Drifter4242 <davehind@yahoo.co.uk> Co-authored-by: Dave Hind <davehind42@gmail.com> Co-authored-by: Valentin Roussellet <vroussellet@apple.com> Co-authored-by: Ken Wang <wangxiaotian2007@gmail.com> Co-authored-by: Spass <spass.ni42+github@gmail.com> Co-authored-by: zsun6 <115504839+zsun6@users.noreply.github.com> Co-authored-by: rishabh <rishabhsaiv@gmail.com> Co-authored-by: Yanzhao Wang <yzwang1996@gmail.com> Co-authored-by: devangpratap <115096812+devangpratap@users.noreply.github.com> Co-authored-by: Jasim Kareem <mj9034812@gmail.com> Co-authored-by: Ryan Choi <54298620+Ryan11c@users.noreply.github.com> Co-authored-by: Ronald Mannak <ronaldmannak@me.com> Co-authored-by: Mateusz Rajski <60450261+mateuuszzzzz@users.noreply.github.com> Co-authored-by: Ishaan Potle <ishaanpotle9@gmail.com> Co-authored-by: Moinuddin Shaik <monu.dev@proton.me> Co-authored-by: Moinuddin Shaik <moinuddinmoin1357@gmail.com> Co-authored-by: shaeantoine <42801108+shaeantoine@users.noreply.github.com> Co-authored-by: Nicoló Angileri <nicoloangileri635@gmail.com> Co-authored-by: Sam Small <sam.small@gmail.com> Co-authored-by: Ishti <45158028+ishtihoss@users.noreply.github.com> Co-authored-by: GokayAI <60583610+gokay-ai@users.noreply.github.com> Co-authored-by: Valeriy Sofin <112272443+sofinvalery@users.noreply.github.com> Co-authored-by: Yash Raj Pandey <55940078+devYRPauli@users.noreply.github.com> Co-authored-by: Kaiyu Shi <skyisno.1@gmail.com> Co-authored-by: crabel99 <crabel@mac.com> Co-authored-by: nav <navneet.dagdiya@gmail.com> Co-authored-by: LongYinan <lynweklm@gmail.com> Co-authored-by: Joseph Heck <j_heck@apple.com> Co-authored-by: deadbydawn101 <sanjosepcrepair@gmail.com> Co-authored-by: deadbydawn <gothravenllm@gmail.com> Co-authored-by: Vijay Kapse <vijayskofficial@gmail.com> Co-authored-by: Long Yixing <longyixing331@gmail.com> Co-authored-by: Aditya Dhulipala <adi@dhulipala.xyz> Co-authored-by: Ishaan Mehta <106708534+ishaanzee@users.noreply.github.com> Co-authored-by: Mattt <mattt@me.com> Co-authored-by: Irakli Salia <65120973+Isalia20@users.noreply.github.com> Co-authored-by: Pasha Khosravi <khosravipasha@users.noreply.github.com> Co-authored-by: Angelos Katharopoulos <a_katharopoulos@apple.com> Co-authored-by: bri-prism <288398250+bri-prism@users.noreply.github.com> Co-authored-by: CodeAndCanvas728 <cedar-elk-01@icloud.com> Co-authored-by: SounLabs <159518213+SounLabs@users.noreply.github.com> Co-authored-by: David Dalcu <david.dalcu@gmail.com> Co-authored-by: Akshay Hitesh Shah <109842662+AkshayShah03@users.noreply.github.com> Co-authored-by: Sash <sash@ela.city> Co-authored-by: tc3oliver <57586652+tc3oliver@users.noreply.github.com> Co-authored-by: Magnus Lundstedt <magnus.lundstedt@infidyne.com> Co-authored-by: jrp2014 <jrp2014@users.noreply.github.com> Co-authored-by: Yi <39162855+SleepAviator@users.noreply.github.com> Co-authored-by: “Ronald <“Ronaldmannak@me.com”> Co-authored-by: Ivan Petrukha <93926277+petrukha-ivan@users.noreply.github.com> Co-authored-by: Haojin Yang <yanghaojin@hotmail.com> Co-authored-by: yanghaojin <haojin.yang@greenbit.ai> Co-authored-by: Marcin Markiewicz <75364832+jmmarkiewicz@users.noreply.github.com> Co-authored-by: Corentin Kerisit <corentin.kerisit@gmail.com> Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com> Co-authored-by: Ambisphaeric <270742616+Ambisphaeric@users.noreply.github.com> Co-authored-by: Sankur Kundu <sankur.kundu.tw@gmail.com> Co-authored-by: begcdn <150344566+begcdn@users.noreply.github.com> Co-authored-by: Core <45273146+IamCooller@users.noreply.github.com> Co-authored-by: David Tivris <git@davidtivr.is>

leak global CommandEncoder -- avoid cuda synchronize on process shutdown

b8e3467

- used in mlx-swift - global CommandEncoder are still alive after exit even if no work is happening - the destructor tries to call cuda api, which causes a fatal error

davidkoski commented Sep 8, 2026

View reviewed changes

zcbenz approved these changes Sep 9, 2026

View reviewed changes

zcbenz merged commit 24c699e into ml-explore:main Sep 9, 2026
29 checks passed

This was referenced Sep 10, 2026

Tagged release carrying mlx#3810 (NAX split-K GEMM dtype fix) without the rest of the v0.32.2 vendor bump? ml-explore/mlx-swift#470

Open

patch for Cuda builds ml-explore/mlx-swift#478

Merged

aaishwarymishra pushed a commit to aaishwarymishra/mlx that referenced this pull request Sep 11, 2026

Leak global CommandEncoder on exit (ml-explore#4480)

528ef42

davidkoski mentioned this pull request Sep 14, 2026

[BUG] threads and streams issues #4506

Closed

inureyes mentioned this pull request Sep 18, 2026

fix(CUDA): Do not throw from ~CudaHandle while the driver is unloading #4531

Closed

zcbenz mentioned this pull request Sep 28, 2026

Leak thread local streams on exit in main thread #4576

Merged

davidtai mentioned this pull request Oct 13, 2026

Fix the CUDA exit abort and the completion-worker busy loop, and add three CUDA throughput changes Layr-Labs/mlx#34

Open

ronaldmannak mentioned this pull request Oct 16, 2026

MLX v0.32.3 with JACCL for the Pico stack PicoMLX/mlx-swift#8

Open

Navigation Menu

Repository navigation

leak global CommandEncoder -- avoid cuda synchronize on process shutdown - #4480

leak global CommandEncoder -- avoid cuda synchronize on process shutdown#4480
zcbenz merged 1 commit into
ml-explore:mainfrom
davidkoski:cuda-shutdown

davidkoski commented Sep 8, 2026

Uh oh!

davidkoski Sep 8, 2026

Uh oh!

davidkoski Sep 8, 2026

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

		static thread_local std::unordered_map<int, CommandEncoder> encoders;
		return encoders;

Conversation

davidkoski commented Sep 8, 2026

Uh oh!

davidkoski Sep 8, 2026

Choose a reason for hiding this comment

Uh oh!

davidkoski Sep 8, 2026

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants