MLX v0.32.0 adds CUDA JIT kernel caching, Metal zero-copy DLPack, and Windows JIT support

MLX v0.32.0, released July 7, 2026, adds CUDA JIT compilation for quantized matmul kernels, Metal zero-copy DLPack sharing, and Windows JIT compiler support, along with a large set of backend fixes.

CUDA and Metal backend changes

The CUDA backend now JIT-compiles qmm_naive, qmm_sm80, qmm_sm90, and gather_gemm kernels, and caches JIT-compiled CUDA kernels. The gemv kernel is also JIT compilable. A new thread_unsafe_stream API was added, and managed memory is disabled on Tegra devices. The qmm_naive kernel got fixes for K-tail dispatch with FP quantized kernels and for scale and bias loads at tile boundaries. qmv now supports global scale, and gather_mm was fixed. CUDA all-reduce planning for large inputs was corrected, and a race in MeshImpl::all_reduce was resolved. A barrier was added to JACCL.

The Metal backend now supports zero-copy DLPack sharing, and a new metal::set_metallib_path() function overrides the metallib search path. Custom Metal kernels can be exported without a Metal backend, and a use-after-free bug was fixed when a custom kernel is called with different dtypes in one graph. Depthwise conv2d handles non-multiple-of-8 spatial dimensions, and the backend is enabled by default on iOS. The NAX kernel name mismatch in gather_qmm was fixed, and int16 overflow in NAX qmm edge-tile bounds was corrected. Tensor-scale nvfp4 is rejected in qqmm.

CPU and core fixes

The CPU allocator without GPU now has buffer caching, synchronized with CPU streams. The no-GPU cache eviction is synchronized with CPU streams. The same thread-local approach used elsewhere was applied to the CPU. Contiguity is computed from the actual occupied data. The compile cache rolls back an entry when the first trace throws, and threaded compile cache cleanup was fixed. Tracing state is now thread-local to fix a data race. The JIT compiler works on Windows, and the prebuilt CPU compile preamble is not used when headers are installed. A missing break in a bool switch case that fell through to uint8_t was fixed. The steel GEMM safe load offset and a typo in steel BaseMMAFrag::load_safe were corrected.

API and correctness updates

New functions include determinant and sign-log-determinant in mlx.core.linalg, flip and unstack, and array API standard aliases for existing ops. The finfo struct now has bits and smallest_normal. Added isdtype, result_type, and can_cast, plus empty, empty_like, astype, and matrix_transpose as pure aliases. mx.asarray gained a copy keyword. DLPack-compatible array imports were improved. The device_count() function returns 0 when there is no GPU. Shape dimension overflow over int32 is detected at MLX compute-shape boundaries, and clearer errors are raised when shape dimension overflows int32. Matvec row offset and gather-MM batch stride int32 overflow were fixed, along with signed-integer overflow in roll and tile shape arithmetic. arange raises on step equal to 0 instead of undefined behavior. The sort VJP transposes the permutation. JVPs for power, divmod, slice_update, select, arctan2, masked_scatter, and bitwise ops were fixed with partly traced inputs. Complex VJPs for log and exp were corrected. Cumulative max and min scans got a VJP. The rope single token multiple sequences bug was fixed, and qvm_split_k incorrect batch stride calculation was corrected. slice update with op indexing bug was fixed. scatter_prod GPU hang on NaN with contention was resolved. The fast layer_norm VJP had an intermittent wrong bias gradient due to a Metal WAR hazard, now fixed. Non-finite float constants emit valid kernel source. GroupNorm and InstanceNorm validate inputs. nll_loss axis parameter was fixed. Adafactor factored update works for params with more than 2 dims. Triplet loss p-norm computation was corrected. SinusoidalPositionalEncoding handles invalid dimensions. Gaussian nll loss default reduction was corrected from none to mean. Optimizer epsilon validation message and AdaDelta eps doc were fixed. ConvTranspose kernel size repr was corrected. The step activation is inclusive at the threshold. tree_map_with_path and tree_merge support namedtuples and tuple subclasses. repeat zero axis shape was fixed. The dimension count in orthogonal initializer error message was corrected. Q4_1 GGUF loading was fixed. The bool switch case missing break was fixed. The MLX_SDPA_BLOCKS env var overrides the 2-pass vector kernel block count. The launch command no longer automatically prepends python, and env variables propagate width. NAX requires setting MACOSX_DEPLOYMENT_TARGET to 26.2. The metallib linking passes the deployment target. The build generates qmm implementations with cmake, reuses the nightly build ccache for release, and uses uv in macOS CI. The actions checkout action was bumped from 6 to 7. The nanobind dependency was updated to 2.13.0.