MLX Updates & Release Notes

What's Changed

• Preserve subnormal float values when casting to bool by @reckylurker in #4224
• Support assigning through a bare Ellipsis index by @Adityaj0 in #4314
• Fix divmod truncating the quotient for floats by @ayaangazali in #4108
• Add force_fused option to scaled_dot_product_attention by @hojin12312 in #4185
• Reject a negative eps in the normalization layers by @ayaangazali in #4312
• Bound GGUF metadata string/array values against the file mapping by @x14ngch3n in #4212
• Read each K/V byte once in gqa-8 decode attention by @dudududukim in #4077
• Fix fft vmap and jvp for transforms over a subset of axes by @kapellirohith in #4138
• Fix median dropping NaN by @devteamaegis in #4146
• Fix the CPU scan over a size one axis with a padded stride by @kapellirohith in #4139
• Validate the optimizer betas at construction by @ayaangazali in #4310
• RMSNormVJP backward writes a full {n_rows, D} gw_temp intermediate by @JasonHonKL in #4293
• [Bug]: add default none value to axis parameter of the take_along_axis by @aaishwarymishra in #4357
• Add a fused full-attention path for head_dim 256 on NAX devices by @wyanzhao in #3842
• Update nanobind to 2.15.0 by @XXXXRT666 in #4337
• Skip unnecessary simdgroup computations for quantised MOE matmuls on NAX by @RohanGautam in #4352
• Add AI usage policy by @zcbenz in #4331
• Raise cpu stream errors from synchronize by @robertomeroni in #4338
• chore: Validate eps in Adam at construction by @vraj00222 in #4361
• Bound winograd conv2d working set by tiling the batch by @Gusanidas in #4102
• Use a 32-row block in qmm_t_nax when one block covers all of M by @dwijenpatel in #4171
• Deduplicate fftshift and ifftshift by @Adityaj0 in #4318
• Fix Log and Equal is_equivalent ignoring primitive state by @kapellirohith in #4266
• Stabilize reduced-precision InstanceNorm by @ternaus in #4230
• Normalize negative axes in sort and argsort by @deBrian07 in #4332
• Clean up main thread compile cache before python interpreter shuts down by @zcbenz in #4373
• chore: Check malformed jaccl hostfile that miss rdma in pairs by @erwinzhang7 in #4284
• Round mxfp8 block scales up to avoid saturation by @dhiltgen in #4353
• Add support for the array_namespace_info by @aaishwarymishra in #4334
• Stop a failed CUDA graph commit from poisoning the encoder by @strayberry in #4356
• Fix quantized kernels in JIT build by @dwijenpatel in #4372
• [Metal][Performance] Avoid zero work in stride-2 ConvTranspose3d by @ternaus in #4343
• [CUDA] Ce fused kernel by @nastya236 in #3947
• Fix cpu exclusive scan for complex numbers by @ayaangazali in #4272
• Support Relocatable CUDA DLLs on Windows by @dhiltgen in #4382
• Use cast_to for fused AsType in compiled Metal kernels by @katlun-lgtm in #4351
• Declare DLPackCompatible protocol members as methods, not settable attributes by @Adityaj0 in #4384
• Fix quantizing sliced arrays by @zcbenz in #4381
• Fix einsum dropping a trailing empty subscript by @Adityaj0 in #4299
• Add script to run python tests by @zcbenz in #4393
• Hold GIL in AttachedData destructor by @zcbenz in #4391

New Contributors

• @hojin12312 made their first contribution in #4185
• @x14ngch3n made their first contribution in #4212
• @vraj00222 made their first contribution in #4361
• @ternaus made their first contribution in #4230
• @strayberry made their first contribution in #4356

Full Changelog: v0.32.1...v0.32.2

Original source

What's Changed

• Fix int64 type cast error when loading GGUF metadata arrays by @danlee2002 in #3823
• Document default values in normalization layer docstrings by @Pablosinyores in #3819
• Warn at configure time when NAX kernels are disabled by @pierre427 in #3824
• [CUDA][Improvement] RMSNorm forward speed up by @nastya236 in #3850
• Fix captured random state in compile by @angeloskath in #3828
• Fix JIT preamble header filter matching project paths containing "Xcode" by @apocryphx in #3873
• Document default value of p in dropout layer docstrings by @ayaangazali in #3870
• [WIP] [CUDA] fsdp by @nastya236 in #3768
• Fix triplet_loss docstring to document the reduced output shape by @ayaangazali in #3884
• Zero-copy CPU import: mx.array(host_buffer, copy=False) on unified memory by @HaoXuAI in #3872
• Round MLX_SDPA_BLOCKS up to a multiple of 32 by @pierre427 in #3875
• Reuse Metal WAR tracking hash tables by @neilmehta24 in #3882
• Use unroll_count(4) for the NAX attention [email protected] loop by @wyanzhao in #3843
• metal: add gemv_wide for fp16/bf16 matmuls of a few rows by @jessegross in #3888
• Fix broken docstring rendering in Linear and RNN by @ayaangazali in #3890
• Fix Adamax betas docstring and MultiOptimizer filters type by @ayaangazali in #3889
• Fix incorrect nvfp4 quantized_matmul through the split-K path by @metascroy in #3854
• [Metal] Avoid regex in custom kernel name generation by @aleroot in #3869
• Fix prod dtype promotion when reducing a size-1 axis by @eyupcanakman in #3898
• [CUDA] Fix grid overflow in gemm conv unfold kernels for >= 65,536 output positions by @AdamDLuz in #3893
• [CUDA] columnwise quantize with tma by @nastya236 in #3157
• metal: reduce NVFP4 scales per 16-lane group by @jessegross in #3934
• Update homebrew in CI by @angeloskath in #3946
• Make index autodiff errors explicitly recommend stop_gradient by @dogukanveziroglu in #3820
• Making JACCL coordinator optional by @angeloskath in #3899
• Fix docstring mismatches in the Python bindings by @ayaangazali in #3948
• Raise a clear error for an invalid quantization mode in nn layers by @ayaangazali in #3914
• Fix incorrect examples and outputs in the usage docs by @ayaangazali in #3956
• Skip test_gather_qmm_sorted cpu test on M1 mac by @zcbenz in #3973
• Fixes an axis mismatch bug in matrix norm for case -1 and 1 by @danlee2002 in #3827
• Fix BatchNorm running variance estimator by @ishtihoss in #3817
• Fix build error caused by TMA macro guard by @zcbenz in #3988
• Fix Glorot/He uniform init docstrings to label the uniform bound, not sigma by @vineethsaivs in #3831
• Fix custom metal kernel cache collision for same name, different source by @katlun-lgtm in #3833
• Fix log_cosh_loss docstring to document the element-wise loss by @winklemad in #3846
• Fix step activation docstring to match >= threshold behavior by @ayaangazali in #3902
• docs: document the reduced-precision float32 default and MLX_ENABLE_TF32 by @stoyoda0012-cyber in #3894
• Fix InstanceNorm Shape docstring to require at least 3 dimensions by @ayaangazali in #3903
• Fix filter_and_map docstring argument order for filter_fn and is_leaf_fn by @ayaangazali in #3906
• Export C++20 requirement to CMake consumers by @PhysicistJohn in #3971
• Fix implicit thread address space qualifier becoming explicit in metal 4.1 by @louen in #3963
• Fix Transformer ignoring a custom encoder or decoder with no parameters by @ayaangazali in #3962
• docs: remove references to removed --no-verify-script launch flag by @latent-9 in #3959
• added eye(0) support and tests by @aaishwarymishra in #3952
• Refactor the JACCL ring and add threads for the multiple rings by @angeloskath in #3900
• Fix shapeless matmul with dynamic batch dimensions by @varshneydevansh in #3813
• Fix JIT build with old macOS SDK by @metascroy in #3853
• Fix cholesky_inv documented argument name to match the binding by @ayaangazali in #3950
• Add new_thread_unsafe_stream to the devices and streams docs by @ayaangazali in #3968
• Add gather_qqmm by @zcbenz in #3757
• Fix missing printoptions doc page from autosummary filename collision by @ayaangazali in #3985
• Document ThreadLocalStream and iinfo in the API reference by @ayaangazali in #3986
• Make "stop" optional in arange by @aaishwarymishra in #3982
• Add softsign to the nn functions docs by @ayaangazali in #3989
• Fix broken all_sum and Group references in the data parallelism example by @ayaangazali in #3996
• Fix ast.metal_kernel typo in the custom Metal kernels guide by @ayaangazali in #3997
• Fix unresolved mx.array docstring references by @ayaangazali in #3990
• python: fix bfloat16 buffer format itemsize mismatch by @reckylurker in #3975
• Template Metal complex scalar lanes by @PhysicistJohn in #3970
• Pad 2D conv input channels to reach the specialized Metal kernel by @eyupcanakman in #3904
• Support head dimension 96 in Metal full attention by @dhiltgen in #3943
• Fix mx.remainder floored-mod for float16/bfloat16 on CPU by @sashko-zakharchuk in #3976
• Fix sorted gather_mm activation row stride by @metascroy in #3960
• Fix state corruption when a primitive throws during eval by @WindChimeRan in #3675
• added complex support by @aaishwarymishra in #3984
• Validate freeze and unfreeze keys against the whole model when recursing by @ayaangazali in #3966
• Fix mlx.launch --python: flag is parsed but never forwarded to the launch script by @jonathan308 in #4002
• Fix broken fully_shard reference in FullyShardedModule docstring by @ayaangazali in #4007
• Use the current interpreter in the comparative benchmark runner by @ayaangazali in #4016
• Document MLX environment variables by @XXXXRT666 in #4000
• Fix CUDA batched GEMV grid overflow by @jasp-nerd in #3929
• Fix all_gather benchmark collapsing its input to a scalar by @ayaangazali in #4017
• Check threadgroup size in the 1-pass sdpa_vector dispatch by @apocryphx in #4018
• Align example projects with the Python 3.10 minimum by @ayaangazali in #4024
• Fix C++ benchmarks failing to build on overloaded astype by @ayaangazali in #4025
• Fix conv_transpose maxBufferLength failures on Metal via tiled unfold by @eyupcanakman in #3845
• Give each host a unique rank in Hostfile.from_list by @ayaangazali in #4027
• Fix crash when reporting partial rings in mlx.distributed_config by @ayaangazali in #4026
• docs: Do not pass MLX_METAL_FAST_SYNCH=1 by default by @katlun-lgtm in #4005
• Fix signed-integer overflow in convolution shape arithmetic by @eyupcanakman in #3938
• Template Metal C2C FFT scalar lanes by @PhysicistJohn in #3969
• Fix segfault in expand_dims for out of bounds negative axes by @Gusanidas in #4021
• Add optional dtype parameter to zeros_like and ones_like by @reckylurker in #4028
• Fix mx.longsumexp output shape issue. by @JasonHonKL in #4030
• Treat backend and envs as optional when parsing a hostfile by @ayaangazali in #4039
• Report bad hostfiles as CLI errors in mlx.launch by @ayaangazali in #4040
• Fix ceil error message to say ceil instead of floor by @PhysicistJohn in #4042
• Fix installed static MLX package on Windows by @yingjiacai in #3848
• Derive the qmv fast path K alignment from bits by @jlee5814 in #3965
• Fix eigh UPLO and zero-size eigh/svd on the CPU by @kapellirohith in #3834
• Add metadata to exported functions by @eyupcanakman in #3855
• Fix empty matrix case in cholesky by @prady0t in #4033
• Reject a negative max_norm in clip_grad_norm by @ayaangazali in #4058
• Build extensions with the interpreter running the build by @ayaangazali in #4057
• Reject negative dimensions in broadcast_to and random shapes by @ayaangazali in #4046
• Report bad host arguments as CLI errors instead of tracebacks by @ayaangazali in #4045
• Fix Metal FFT for sizes above 2**20 by @adityasingh2400 in #4013
• Request MSL 4.1 from the runtime compiler on macOS 27 by @magnuslundstedt in #4052
• Fix concurrent Metal kernel cache lookup by @kitty-eu-org in #4043
• Refactor CPU ArgReduce dtype dispatch by @PhysicistJohn in #4066
• Refactor NumberOfElements dtype dispatch by @PhysicistJohn in #4064
• Refactor CPU Select dtype dispatch by @PhysicistJohn in #4065
• Raise qmv batch limit for large matrices on M5-class GPUs by @pierre427 in #3791
• Enable half-precision complex Metal FFT kernels by @PhysicistJohn in #3981
• Fix: mx.distributed.sum_scatter crashes on a scalar(0-d) by @JasonHonKL in #4071
• Return an empty result from pinv for zero-size inputs by @ayaangazali in #4069
• Handle zero-size inputs in eig and qr on the CPU by @ayaangazali in #4068
• Normalize biases before encoding in gather_qmm_rhs by @dudududukim in #4056
• Keep randint samples inside [low, high) by @adityasingh2400 in https://github.com//pull/4012
• Allow saving empty arrays to npy and safetensors by @ayaangazali in #4080
• Fix DeviceType annotations for device and stream arguments by @XXXXRT666 in #4059
• Refactor CPU DivMod dtype dispatch by @PhysicistJohn in #4088
• Refactor CPU SegmentedMM dtype dispatch by @PhysicistJohn in #4087
• Refactor CPU Arange dtype dispatch by @PhysicistJohn in #4085
• Fix Device and Stream lexicographic ordering by @wgu9 in #4086
• Enable complex64 scatter addition on GPU by @PhysicistJohn in #4078
• Optimize large NVFP4 QMV on M5 Max by @dhiltgen in #3961
• ring: fail on peer disconnect instead of hanging forever by @erwinzhang7 in #4060
• Fix: propagate NaN when dealing with cummax and curmin by @JasonHonKL in #4047
• Refactor CPU equal_nan dtype dispatch by @PhysicistJohn in #4095
• Refactor CPU Matmul dtype dispatch by @PhysicistJohn in #4103
• Refactor CPU Gather dtype dispatch by @PhysicistJohn in #4105
• Refactor generic CPU unary dtype dispatch by @PhysicistJohn in #4106
• Refactor generic CPU binary dtype dispatch by @PhysicistJohn in #4107
• Validate pooling kernel size, stride and padding by @ayaangazali in #4101
• Only reject max and min when a reduced axis is empty by @ayaangazali in #4079
• Fix SinusoidalPositionalEncoding silently overriding explicit scale=0.0 by @YassineMA03 in #4098
• Fix any and all treating -0.0 as nonzero by @ayaangazali in #4090
• Fix Eigh CPU dtype error message by @PhysicistJohn in #4130
• Fix random distribution dtype error messages by @PhysicistJohn in #4131
• Fix safetensors dtype error formatting by @PhysicistJohn in #4132
• Fix no-GPU synchronize error formatting by @PhysicistJohn in #4133
• Fix: mx.isinf doens't consider imaginary plane. by @JasonHonKL in #4092
• Refactor CPU Copy destination dtype dispatch by @PhysicistJohn in #4104
• Refactor CPU GatherAxis dtype dispatch by @PhysicistJohn in #4109
• Refactor inexact CPU unary dtype dispatch by @PhysicistJohn in #4111
• Refactor CPU comparison dtype dispatch by @PhysicistJohn in #4112
• Replace Ring's local dtype macro with common dispatch by @PhysicistJohn in #4120
• Refactor CPU ArgSort dtype dispatch by @PhysicistJohn in #4123
• Refactor CPU Softmax dtype dispatch by @PhysicistJohn in #4127
• Use common dtype dispatch for Metal Arange by @PhysicistJohn in #4128
• Pick BM from rows per expert in gather_qmm_rhs_nax by @dwijenpatel in #4023
• Enable complex64 take_along_axis backward on Metal by @PhysicistJohn in #4094
• Add mx.searchsorted with CPU, Metal and CUDA kernels by @erwinzhang7 in #4035
• Stop integer power from hanging on a negative exponent by @ayaangazali in #4100
• Fix second order gradients for sort, partition, topk and cummax/cummin by @Adityaj0 in #4117
• Fix segfault on negative out of bounds axes in take_along_axis/put_along_axis by @Adityaj0 in #4118
• Deduplicate slow CPU convolution dtype dispatch by @PhysicistJohn in #4126
• Use common dispatch for compiled constants by @PhysicistJohn in #4129
• Fix set_printoptions example by @PhysicistJohn in #4140
• Clarify export output type diagnostic by @PhysicistJohn in #4141
• Fix integer keys collision in tree_unflatten by @Solaris-star in #3878
• Fix fence tracking for donated dynamic slice offsets by @michaelellis003 in #4099
• Deduplicate distributed all-reduce construction by @PhysicistJohn in #4144
• Deduplicate distributed socket errors by @PhysicistJohn in #4145
• Remove redundant Metal bitwise dispatch by @PhysicistJohn in #4147
• Deduplicate custom Metal kernel size validation by @PhysicistJohn in #4149
• Deduplicate CUBLAS compute type cases by @PhysicistJohn in #4151
• Refactor CPU Scatter dtype dispatch by @PhysicistJohn in #4155
• fix(cpu): name the right routine in LUF and SVD error messages by @devteamaegis in #4167
• Refactor CPU ArgPartition dtype dispatch by @PhysicistJohn in #4156
• add complex to scalar to avoid raising warnings by linters by @aaishwarymishra in #4168
• Refactor CPU Copy source dtype dispatch by @PhysicistJohn in #4170
• Deduplicate Metal scan encoder setup by @PhysicistJohn in #4154
• [Metal] Skip empty NAX GEMM output groups by @XXXXRT666 in #3941
• docs: fix stale CLI invocations in LLM inference example by @deBrian07 in #4172
• Increase ccache size in CI by @zcbenz in #3999
• Fix hadamard_transform on the GPU for n = m with no power-of-2 factor by @axiom-of-choice in #4054
• Refactor CPU Partition dtype dispatch by @PhysicistJohn in #4175
• Refactor CPU ScatterAxis and MaskedScatter dtype dispatch by @PhysicistJohn in #4176
• docs: document softmax's precise argument by @apocryphx in #4178
• Send out of range trig arguments to libm by @ayaangazali in #4157
• Add reflect and symmetric padding modes to mx.pad by @katlun-lgtm in #3608
• docs: add a usage page on growing arrays by @apocryphx in #4019
• Fix and enable non-transposed NAX qmm by @gordofreemo in #4051
• Add scatter reduce for JACCL by @angeloskath in #3901
• Fix diag for zero-size input by @devteamaegis in #4165
• Add an inverse-CDF path to categorical sampling by @deescondrillas in #4177
• Fail when a requested RDMA device is not found by @erwinzhang7 in #4180
• Bound GGUF tensor data offsets against the file mapping by @robertomeroni in #4179
• Shift by the max in log_softmax by @ayaangazali in #4169
• Shift by the max in cross_entropy by @ayaangazali in #4188
• Fix linalg.norm(x, keepdims=True) not keeping dims by @devteamaegis in #4166
• Refactor inexact CPU binary dtype dispatch by @PhysicistJohn in #4194
• Implement batched matmul for large 1D dot products by @Ved235 in #3580
• Re-disable qmm_n_nax and fix group_size < 64 by @PhilipJohnBasile in #4202
• fix: shell-quote the working directory in the distributed launch script by @bunlongheng in #4189
• Report when no usable GID is found by @erwinzhang7 in #4191
• Keep double precision for python floats in float64 operations by @nileshpatil6 in #4173
• Jaccl ring multi wire fix by @angeloskath in #4193
• Report why creating a queue pair failed by @erwinzhang7 in #4209
• Add antialias support to nn.Upsample for linear and cubic modes by @lyonsno in #3677
• Require pytorch 2.12 for metal dlpack tests by @zcbenz in #4206
• Remove std::optional in C++ public interface by @zcbenz in #4207
• Fix SliceUpdate JVPs with one traced input by @PhysicistJohn in #4200
• Raise instead of hanging when reflect/symmetric pad gets an empty axis by @codeAnqiang-ma in #4223
• Reuse Stockham twiddles in large batched Bluestein FFTs by @PhysicistJohn in #4084
• Validate decay_steps and step_size in the schedulers by @ayaangazali in #4217
• Add runnable distributed examples by @erwinzhang7 in #4219
• Break monolithic MTLResidencySet into smaller sets by @JamesMcCarthy44 in #4211
• [CUDA][Improvement] RMSNorm backward by @nastya236 in #3881
• Fix MultiOptimizer on models containing empty modules by @erenmenges in #4215
• Return tuple in meshgrid by @JasonHonKL in #4229
• Add endpoint parameter to linspace by @AKnassa in #4184
• Fix vmap of partition/argpartition dropping the kth argument by @Adityaj0 in #4116
• Fix nan_to_num replacing inf with 0 for float16 and bfloat16 by @codeAnqiang-ma in #4222
• Fix einsum not broadcasting batch dimensions in batched tensordot by @Adityaj0 in #4125
• Dequantize in float32 by @RohanGautam in #4241
• Reject complex in erf and erfinv by @ayaangazali in #4243
• Make abs a no-op for unsigned types in the simd helpers by @ayaangazali in #4240
• Fix quantize matrix multiplication floor issue by @JasonHonKL in #4251
• Only use MPI backend for world size > 1 by @erwinzhang7 in #4210
• Reject complex in expm1, sigmoid and arctan2 by @ayaangazali in #4257
• Decompose small kernel-depth 3D convs into 2D convs by @katlun-lgtm in #3785
• Fix Metal sort of a view with a negative stride by @robertomeroni in #4252
• Mirror the depth axis in the decomposed 3D conv when flipped by @ayaangazali in #4277
• Fix Metal row reductions on negative-stride views by @FU-max-boop in #4267
• [CUDA] Fix custom kernel cache collision for same name, different source by @HaoXuAI in #4273
• Fix ops rejecting integers larger than INT32_MAX by @FeliGame in #4255
• Fix var/std for complex numbers by @ayaangazali in #4260
• Fix int32 overflow in conv padded input and pad shapes by @eyupcanakman in #4258
• Reject complex in remainder by @ayaangazali in #4270
• Compare the macOS SDK version as a version when gating JACCL by @erwinzhang7 in #4286
• Clamp ring socket transfers so a payload of 2 GiB or more can be sent by @erwinzhang7 in #4281
• Replace hand-rolled axis normalization with normalize_axis_index in split/unstack/partition/topk by @Adityaj0 in #4288
• Remove grouped output in CI by @zcbenz in #4195
• [CUDA] Fix finding cuda 13 headers in JIT compilation by @zcbenz in #3995
• Refactor wheel building script by @zcbenz in #3818
• Make mx.compile cache erasing thread safe by @zcbenz in #4248
• Add builds for free-threaded python by @zcbenz in #3812
• Fix int32 overflow in concatenate/repeat/kron by @ayaangazali in #4303
• Widen list elements that do not fit in int32 to int64 by @ayaangazali in #4305
• Propagate CPU errors to events by @zcbenz in #3742
• Fix mx.arange dtype inference overflow regression by @shoutoutuoadi325 in #4324
• Add workflow to update pull request limit bypass list by @zcbenz in #4320
• Support head dimension 72 in Metal full attention by @dhiltgen in #4330

New Contributors

• @pierre427 made their first contribution in #3824
• @apocryphx made their first contribution in #3873
• @ayaangazali made their first contribution in #3870
• @HaoXuAI made their first contribution in #3872
• @neilmehta24 made their first contribution in #3882
• @wyanzhao made their first contribution in #3843
• @metascroy made their first contribution in #3854
• @aleroot made their first contribution in #3869
• @AdamDLuz made their first contribution in #3893
• @vineethsaivs made their first contribution in #3831
• @winklemad made their first contribution in #3846
• @stoyoda0012-cyber made their first contribution in #3894
• @PhysicistJohn made their first contribution in #3971
• @latent-9 made their first contribution in #3959
• @aaishwarymishra made their first contribution in #3952
• @reckylurker made their first contribution in #3975
• @sashko-zakharchuk made their first contribution in #3976
• @WindChimeRan made their first contribution in #3675
• @jonathan308 made their first contribution in #4002
• @jasp-nerd made their first contribution in #3929
• @Gusanidas made their first contribution in #4021
• @JasonHonKL made their first contribution in #4030
• @yingjiacai made their first contribution in #3848
• @jlee5814 made their first contribution in #3965
• @kapellirohith made their first contribution in #3834
• @prady0t made their first contribution in #4033
• @magnuslundstedt made their first contribution in #4052
• @kitty-eu-org made their first contribution in #4043
• @dudududukim made their first contribution in #4056
• @erwinzhang7 made their first contribution in #4060
• @YassineMA03 made their first contribution in #4098
• @dwijenpatel made their first contribution in #4023
• @Adityaj0 made their first contribution in #4117
• @Solaris-star made their first contribution in #3878
• @michaelellis003 made their first contribution in #4099
• @devteamaegis made their first contribution in #4167
• @deBrian07 made their first contribution in #4172
• @axiom-of-choice made their first contribution in #4054
• @gordofreemo made their first contribution in #4051
• @deescondrillas made their first contribution in #4177
• @robertomeroni made their first contribution in #4179
• @Ved235 made their first contribution in #3580
• @PhilipJohnBasile made their first contribution in #4202
• @bunlongheng made their first contribution in #4189
• @nileshpatil6 made their first contribution in #4173
• @codeAnqiang-ma made their first contribution in #4223
• @JamesMcCarthy44 made their first contribution in #4211
• @erenmenges made their first contribution in #4215
• @AKnassa made their first contribution in #4184
• @RohanGautam made their first contribution in #4241
• @FU-max-boop made their first contribution in #4267
• @FeliGame made their first contribution in #4255
• @shoutoutuoadi325 made their first contribution in #4324

Full Changelog: v0.32.0...v0.32.1

Original source

What's Changed

Generate qmm implementaions with cmake by @zcbenz in #3424
Enable swap for all CI building CUDA by @zcbenz in #3437
Bump minor by @angeloskath in #3438
[CUDA] Fix qmm_naive K-tail dispatch for FP quantized kernels by @Lyxot in #3445
Keep gguflib input-validation asserts active in release builds by @qflen in #3436
Reuse nightly build's ccache for release by @zcbenz in #3458
Add barrier to JACCL by @Isalia20 in #3459
Add determinant and sign-log-determinant functions to mlx.core.linalg by @abhilashreddys in #3416
Define ST_F8_E8M0 by @pcuenca in #3448
Clearer error when shape dimension overflows int32 by @serenposh in #3425
[CUDA] Fix half type matmul in cutlass kernels by @zcbenz in #3469
Fix indexing bug in slice update with op by @angeloskath in #3483
Make device_count() return 0 when there is no GPU by @zcbenz in #3486
Compute contiguity from the actual occupied data by @sofinvalery in #3475
Do not use prebuilt cpu compile preamble when headers are installed by @zcbenz in #3463
[CUDA] Separate main loop into a function in qmm by @zcbenz in #3443
test: Upcast the random numbers before computing their average by @sofinvalery in #3488
Pass deployment target when linking metallib by @dhiltgen in #3501
Fix rope single token multiple sequences by @angeloskath in #3498
Fix qvm_split_k incorrect batch stride calculation by @angeloskath in #3497
Fix scatter_prod GPU hang on NaN with contention by @tillahoffmann in #3492
[jaccl] Fix race on local_staging in MeshImpl::all_reduce by @kernelpool in #3451
Add MLX_SDPA_BLOCKS env var for 2-pass vector kernel block-count override by @adurham in #3455
mlx launch clean by @nastya236 in #3513
[CUDA] Fix gather_mm by @zcbenz in #3503
[CUDA] Guard qmm_naive scale and bias loads at tile boundaries by @Lyxot in #3509
win: fix cuda build by @dhiltgen in #3532
Improve DLPack-compatible array imports by @XXXXRT666 in #3495
Apply the same thread-local approach to the CPU by @angeloskath in #3537
removed automatic prepending python in mlx launch by @nastya236 in #3536
ci: Show stack trace on crash by @zcbenz in #3538
Handle non-multiple-of-8 spatial dims in depthwise conv2d Metal path by @qflen in #3446
Remove reference to groups in Conv3D extra_repr by @GriffinMB in #3559
Add buffer caching to no_gpu CPU allocator by @dhiltgen in #3554
Fix steel GEMM safe load offset by @sofinvalery in #3560
Fix singleton lifetime issues at process exit by @dhiltgen in #3555
Fix doubled-word typos in docstrings by @LeSingh1 in #3561
Fix off_x/off_y typo in steel BaseMMAFrag::load_safe by @mdegans in #3565
Synchronize no-GPU cache eviction with CPU streams by @dhiltgen in #3566
Add JIT compiler support for Windows by @dhiltgen in #3556
[Metal] Reject tensor-scale nvfp4 in qqmm by @Brooooooklyn in #3551
Detect int32 shape-product overflow at MLX compute-shape boundaries by @qflen in #3524
Add copy keyword to mx.asarray by @eyupcanakman in #3510
Fix typos in backend code comments by @adityasingh2400 in #3581
Implement output_shapes for GatherMM and GatherQMM by @dexwritescode in #3485
Fix activation docstring math in log_softmax and gelu_approx by @adityasingh2400 in #3590
docs: expose missing Python APIs by @XXXXRT666 in #3598
Fix CUDA all-reduce planning for large inputs by @sofinvalery in #3603
Use uv in macOS CI by @zcbenz in #3491
Fix HDRS_LIST building when paths in header inclusion tree include spaces by @lancelotblanchard in #3607
Enable the Metal backend by default on iOS by @eyupcanakman in #3617
Fix int32 overflow in matvec row offset and gather-MM batch stride by @aicayzer in #3609
Fix signed-integer overflow (UB) in roll and tile shape arithmetic by @devYRPauli in #3604
Fix complex VJPs for log and exp by @CameronChurchwell in #3605
added env variables to propagate width by @nastya236 in #3618
Correct p-norm computation in triplet_loss by @pchintar in #3613
Fix threaded compile cache cleanup by @lucasnewman in #3628
Handle invalid dimensions in SinusoidalPositionalEncoding by @pchintar in #3615
Correct grammatical typo in label smoothing ValueError by @madhav1k in #3641
Roll back compile cache entry when the first trace throws by @adityasingh2400 in #3635
Raise on arange with step == 0 instead of undefined behavior by @devYRPauli in #3640
Fix data race in tracing state by making it thread-local by @aicayzer in #3638
Add missing break in bool_ switch case that would result in falling through to the uint8_t case by @psolanki in #3655
Fix axis param in nll_loss by @ishtihoss in #3651
Fix Adafactor factored update for params with more than 2 dims by @ishtihoss in #3652
Emit valid kernel source for non-finite float constants by @tillahoffmann in #3648
Validate inputs in GroupNorm and InstanceNorm by @ishtihoss in #3653
Fix flaky test_siblings_without_eval by @zcbenz in #3621
Fix intermittent wrong bias gradient in fast::layer_norm VJP (Metal WAR hazard) by @tillahoffmann in #3630
NAX requires setting MACOSX_DEPLOYMENT_TARGET=26.2 by @zcbenz in #3622
Catch error in CommandBuffer and poison the events by @zcbenz in #3523
[Improvement] CopyType::Vector in concatenate if axis=0 and contiguous by @nastya236 in #3663
Fix Q4_1 GGUF loading by @ricky-chaoju in #3664
Fix tree_map_with_path for namedtuples by @chrismicah in #3674
Fix gather_qmm NAX kernel name mismatch by @scyyh11 in #3632
Add bits and smallest_normal to finfo by @katlun-lgtm in #3679
Fix use-after-free when a custom Metal kernel is called with different dtypes in one graph by @discobot in #3662
Fix dimension count in orthogonal initializer error message by @Pablosinyores in #3693
fix grid for large uncontiguous input by @nastya236 in #3666
Ignore numpy matmul warnings in test_blas by @zcbenz in #3657
Fix repeat zero axis shape by @chrismicah in #3698
Cache JIT-compiled CUDA kernels by @zcbenz in #3587
[CUDA] JIT-compile qmm_naive by @zcbenz in #3576
Support namedtuple and tuple subclasses in tree_merge by @Pablosinyores in #3703
Fix mx.sort vjp to transpose the permutation by @obchain in #3700
Add isdtype, result_type, and can_cast by @katlun-lgtm in #3681
[fix] allow exporting custom Metal kernels without a Metal backend by @xthomaswang in #3650
Make step activation inclusive at the threshold by @Pablosinyores in #3694
Add new_thread_unsafe_stream API by @zcbenz in #3578
Fix JVPs of power, divmod, and slice_update with partly traced inputs by @qflen in #3636
Add flip and unstack by @katlun-lgtm in #3683
Fix int16 overflow in NAX qmm edge-tile bounds by @scyyh11 in #3631
Make gemv JIT compilable by @angeloskath in #3705
Fix gaussian_nll_loss documented default reduction (none -> mean) by @Pablosinyores in #3715
Fix optimizer epsilon validation message and AdaDelta eps doc by @Pablosinyores in #3712
Add array API standard aliases for existing ops by @katlun-lgtm in #3678
Fix ConvTranspose kernel size repr by @ishtihoss in #3724
[CUDA] Disable managed memory on Tegra by @gmondada in #3701
Update nanobind to 2.13.0 by @XXXXRT666 in #3722
Fix JVPs of select, arctan2, masked_scatter, and bitwise ops by @qflen in #3633
array API: add empty, empty_like (pure aliases), astype, matrix_transpose by @katlun-lgtm in #3729
Add metal::set_metallib_path() to override the metallib search path by @olilarkin in #3597
Add Metal DLPack zero-copy sharing by @XXXXRT666 in #3531
Bump actions/checkout from 6 to 7 by @dependabot[bot] in #3741
Fix shapeless compile eliding reductions over size-1 dimensions by @discobot in #3672
[CUDA] JIT-compile qmm_sm80/qmm_sm90/gather_gemm by @zcbenz in #3706
[CUDA] Make qmv support global scale by @zcbenz in #3723
Add VJP for cumulative max and min scans by @devin-lai in #3599
Fix abs jvp for complex inputs by @obchain in #3745
Fix update_modules crash on over-length module lists by @Pablosinyores in #3743
Add matrix_norm to array API namespace by @wgu9 in #3749
Guard scatter_axis against 64-bit outputs on the GPU by @obchain in #3695
Fix build on Apple clang 15 in cpu_nd_array_to_mlx by @obchain in #3751
metal: cache x in registers in rms_single_row to avoid redundant global read by @will-march in #3754
Document row-wise sparsity in nn.init.sparse by @Pablosinyores in #3752
[CUDA] [Improvement] Rope without copy by @nastya236 in #3704
Add small-batch quantized matvec kernel (qmv_wide) by @jessegross in #3764
Fix docstring nits by @ishtihoss in #3758
Fix logsumexp jvp to reduce along the axis by @obchain in #3708
Add vecdot to array API namespace by @wgu9 in #3748
Fix CPU dynamic slice copy bound for collapsed shapes by @Lyxot in #3739
Enable fused SDPA vector kernel for asymmetric Q/V head dims (192, 128) by @yohann-bearzi in #3637
array API: add positive, logical_xor, trunc, count_nonzero, diff, full_like by @katlun-lgtm in #3730
array API: add cumulative_sum and cumulative_prod by @katlun-lgtm in #3731
Fix complex vjps for several unary ops by @obchain in #3766
Fix Upsample align_corners singleton output by @ishtihoss in #3769
Add CI for Windows CUDA build by @zcbenz in #3775
Fix multi-wire recv prefill deadlock in jaccl ring backend by @jasonpaulso in #3654
fix: Quote hostname in mlx.launch ssh commands by @varshneydevansh in #3783
Fix BatchNorm docstring to cover 4D input and test the NHWC path by @Pablosinyores in #3782
Fix kl_div_loss docstring formula and two doc nits by @Pablosinyores in #3788
Fix HardShrink to accept its documented lambd argument by @Pablosinyores in #3786
Refactor macOS CI by @zcbenz in #3806
Bypass the PyTorch index TLS issues by @angeloskath in #3808
Fix fp quantized matvec for output dim < 8 by @katlun-lgtm in #3804
Fix CUDA RMSNorm small-row dispatch by @dhiltgen in #3792
Fix Conv3d docstring argument order (padding before dilation) by @Pablosinyores in #3793
Add math mode option for custom Metal kernels by @shubhxho in #3728
Fix binary_cross_entropy docstring argument order (weights before with_logits) by @Pablosinyores in #3807
Fix compiled kernel correctness for negative-strided inputs by @lyonsno in #3720
Use large runner for Windows CUDA builds by @zcbenz in #3809
Fix wrong type parameter passed to gemm_splitk_nax by @zcbenz in #3810
Document default value of dilation in conv layer docstrings by @Pablosinyores in #3811
Fix CPU gather transposing column-contiguous slices by @tillahoffmann in #3647
Fix infinite norm negative axis mismatch bug for matrices with 2 or more dimensions. by @danlee2002 in #3756
Fix conv2 gradients in grouped strided case on Metal by @ericphanson in #3800

New Contributors

@qflen made their first contribution in #3436
@Isalia20 made their first contribution in #3459
@abhilashreddys made their first contribution in #3416
@serenposh made their first contribution in #3425
@sofinvalery made their first contribution in #3475
@tillahoffmann made their first contribution in #3492
@adurham made their first contribution in #3455
@GriffinMB made their first contribution in #3559
@LeSingh1 made their first contribution in #3561
@mdegans made their first contribution in #3565
@eyupcanakman made their first contribution in #3510
@adityasingh2400 made their first contribution in #3581
@dexwritescode made their first contribution in #3485
@lancelotblanchard made their first contribution in #3607
@aicayzer made their first contribution in #3609
@devYRPauli made their first contribution in #3604
@pchintar made their first contribution in #3613
@madhav1k made their first contribution in #3641
@psolanki made their first contribution in #3655
@ishtihoss made their first contribution in #3651
@ricky-chaoju made their first contribution in #3664
@chrismicah made their first contribution in #3674
@scyyh11 made their first contribution in #3632
@katlun-lgtm made their first contribution in #3679
@discobot made their first contribution in #3662
@Pablosinyores made their first contribution in #3693
@obchain made their first contribution in #3700
@xthomaswang made their first contribution in #3650
@gmondada made their first contribution in #3701
@olilarkin made their first contribution in #3597
@devin-lai made their first contribution in #3599
@wgu9 made their first contribution in #3749
@will-march made their first contribution in #3754
@yohann-bearzi made their first contribution in #3637
@jasonpaulso made their first contribution in #3654
@varshneydevansh made their first contribution in #3783
@shubhxho made their first contribution in #3728
@lyonsno made their first contribution in #3720
@danlee2002 made their first contribution in #3756
@ericphanson made their first contribution in #3800

Full Changelog: v0.31.2...v0.32.0

Original source

Highlights

Wider support for cuda quantized matmuls (#3352, #3268, #3321, #3417, #3255)

MLX can be used by multiple threads for independent computations (#3405, #3348, #3281, #3423)

Added CUDA FFT support

JACCL is now a standalone lib (#3412)

What's Changed

Bump by @angeloskath in #3244

win: re-enable and fix cuDNN performance by @dhiltgen in #3242

Fix crashes in multi-threaded process teardown by @louen in #3167

[CUDA] Add FFT support by @lucasnewman in #3243

[CUDA] Implement MaskedScatter by @Lyxot in #3151

docs: fix PyTorch to MLX conversion example by @LxYuan0420 in #3265

update requirements for Macbook Neo by @tosh in #3257

fix comparison op JVP returning bool tangents instead of input dtype by @mm65x in #3253

fix nn.GRU skipping bhn bias when hidden is None by @mm65x in #3252

[CUDA] Pipelined QMM by @zcbenz in #3255

tests: harden memory leak check in test_siblings_without_eval by @booxter in #3088

Slice update with operation by @angeloskath in #3266

Nax Refactor by @jagrit06 in #3271

Fix building with CUDA toolkit 13.2 by @zcbenz in #3273

[CUDA] fp and int4 quants for qmm_sm80 by @zcbenz in #3268

Fix repr of conv layers by @angeloskath in #3275

Merge DeviceStream into CommandEncoder by @zcbenz in #3264

[CUDA] Search system-installed CUDA toolkit for headers by @zcbenz in #3277

Create default random key lazily by @zcbenz in #3278

Support indexing with any type which implmented index by @aisk in #3210

Fix sort NaN handling for float16 and bfloat16 by @Lyxot in #3269

Use thread local storage for frontend compile cache by @zcbenz in #3280

[Metal][Performance]: Add split-K for quantized matmul (small M) by @Ziqiao-git in #3120

[Metal] Fix depthwise conv 1D kernel name for large variant by @Brooooooklyn in #3289

Fix stale transform copy-chain leaks by @Brooooooklyn in #3290

Implement Pad::vmap to replace NYI stub by @Aristide021 in #3304

logo files by @andresy in #3308

Fix vmap + floor_divide: preserve integer dtype by @robert-johansson in #3292

Fix moved-from shape bug in broadcast_arrays causing vmap bus error by @Aristide021 in #3310

Use nb::ndarray for checking arrays by @zcbenz in #3283

Add output_shapes for AddMM by @pHequals7 in #3262

Manage Metal objects with smart pointers by @zcbenz in #3282

[CUDA] support sorting complex numbers by @Lyxot in #3286

Add norm parameter to FFT transforms (backward/ortho/forward) by @Aristide021 in #3287

Make each thread have its own default stream by @zcbenz in #3281

[CUDA] Implement BlockMaskedMM by @Lyxot in #3299

Fix np bfloat16 misinterpreted as complex by @kellen-sun in #3146

Remove no longer needed const_cast by @zcbenz in #3325

Bump actions/deploy-pages from 4 to 5 by @dependabot[bot] in #3334

Fix use after move by @angeloskath in #3343

Decouple CommandEncoder from Device by @zcbenz in #3316

Add vmap for BroadcastAxes by @angeloskath in #3344

Add fftfreq, rfftfreq and scalar axes for fftshift/ifftshift by @declanhealy2 in #3298

[Metal] Support sorting complex numbers by @Lyxot in #3314

[CUDA] Fallback QMM by @zcbenz in #3315

Make CommandEncoder thread local by @zcbenz in #3348

[CUDA] 3/5/6-bit quants for qmm_naive by @zcbenz in #3352

Fix regression in array creation by @angeloskath in #3353

Use metal as the front-end for the metal linker by @louen in #3354

Add printoptions by @ChristophePRAT in #3333

Add a convenience for making local streams in python by @angeloskath in #3355

Fix CMake finding wrong Python during pip install by @fijimunkii in #3375

[CUDA] Add GatherQMM for quantized gather matmul by @Lyxot in #3321

fix: fail build when Metal compiler header resolution fails by @dogukanveziroglu in #3332

Fix: Correct cross-attention query routing in Post-LN TransformerDecoderLayer by @suryawanshishantanu6 in #3382

[CUDA] Thread safety by @zcbenz in #3367

Fix test "test get streams" missing initialization by @dseredkin in #3376

Conjugate VJP and JVP support by @CameronChurchwell in #3386

Fix int16 overflow in SDPA NAX mask indexing for KV sequences > 32K by @Clydingus in #3361

Avoid joining threads on exit by @zcbenz in #3388

Add clear_streams API for cleanup before exit by @zcbenz in #3395

Update nanobind version to v2.12.0 by @jrp2014 in #3396

Jaccl refactor by @angeloskath in #3412

Fixes for CUDA CI by @zcbenz in #3413

Validate safetensors data offsets by @MillaFleurs in #3364

Validate safetensors data offsets against file boundaries by @matinsaurralde in #3410

Document sort stability and NaN handling by @NeuralNoble in #3400

ThreadLocalStream in C++ by @zcbenz in #3405

Fix jaccl init bug by @angeloskath in #3418

Segmented mm nax kernel by @angeloskath in #3419

[CUDA] gather_mm by @zcbenz in #3414

[CUDA] GatherQMM matrix-matrix sm80/naive path by @Lyxot in #3417

[CUDA] Handle residue k in qmm_naive by @zcbenz in #3379

Speed up NAX split-K by better tuning and routing and fix NAX addmm by @angeloskath in #3422

Make Scheduler::enqueue thread safe by @zcbenz in #3423

Fix flaky TestVmap.test_vmap_masked_scatter by @zcbenz in #3421

Fix synchronize for ThreadLocalStream by @angeloskath in #3429

Fix bytes_per_key truncation in random kernels (Metal + CUDA) by @dogukanveziroglu in #3432

Throw meaningful error when Metal device is not found by @dogukanveziroglu in #3428

Fix kernel cache collision in Compiled constructor by @dogukanveziroglu in #3427

Fix mx.prod vjp for complex types by @CameronChurchwell in #3433

New Contributors

@LxYuan0420 made their first contribution in #3265

@tosh made their first contribution in #3257

@mm65x made their first contribution in #3253

@booxter made their first contribution in #3088

@Ziqiao-git made their first contribution in #3120

@Brooooooklyn made their first contribution in #3289

@Aristide021 made their first contribution in #3304

@pHequals7 made their first contribution in #3262

@declanhealy2 made their first contribution in #3298

@fijimunkii made their first contribution in #3375

@dogukanveziroglu made their first contribution in #3332

@suryawanshishantanu6 made their first contribution in #3382

@dseredkin made their first contribution in #3376

@CameronChurchwell made their first contribution in #3386

@Clydingus made their first contribution in #3361

@jrp2014 made their first contribution in #3396

@matinsaurralde made their first contribution in #3410

@NeuralNoble made their first contribution in #3400

Full Changelog: v0.31.1...v0.31.2

Original source

What's Changed

Bump the patch version by @angeloskath in #3185
Skip Hopper-only kernels in CI by @zcbenz in #3184
[CUDA] Fsdp (easy) by @nastya236 in #3130
Fix ref leak in mx.save/load with file like object by @aisk in #3187
Fix/missing libs in docs by @ChristophePRAT in #3190
feat: adding the bartlett function by @Vlor999 in #3155
Bump actions/download-artifact from 7 to 8 by @dependabot[bot] in #3189
Bump actions/upload-artifact from 6 to 7 by @dependabot[bot] in #3188
[CUDA] Quantized GEMV by @zcbenz in #3180
[CUDA] Use fp16 accumulation for 4-bit quant in GEMV by @zcbenz in #3197
[CUDA] implement Hadamard transform by @Lyxot in #3179
Improve mlx.distributed_config by @angeloskath in #3199
PR#3226 Fix by @MillaFleurs in #3227
PR #3220 LayerNorm VJP returns zeros_like(weight) instead of zeros_like(bias placeholder) by @MillaFleurs in #3231
[CUDA] Faster compilation and batch support in QMV by @zcbenz in #3213
Validate num_splits in split by @MillaFleurs in #3234
Fix return value in einsum_path for simple contractions by @MillaFleurs in #3232
Validate dims in rope by @MillaFleurs in #3230
Fix assigning bool to float16/bfloat16 by @MillaFleurs in #3229
Fix load_weights with strict=False to filter extra weights before update by @gmin7 in #3214
Remove custom fp4/fp8 classes by @zcbenz in #3212
[CUDA] Support 3/5/6-bit quants in QMV by @zcbenz in #3236
Hybrid sharding by @nastya236 in #3194
win: fix cuda build by @dhiltgen in #3204
Remove quantized_utils.cuh by @zcbenz in #3237
[CUDA] Implement SegmentedMM by @Lyxot in #3238
Add initial tuning for M5 pro and max by @jagrit06 in #3211
[CUDA] Use qmv kernel for fp quantizations by @zcbenz in #3239

New Contributors

@ChristophePRAT made their first contribution in #3190
@Lyxot made their first contribution in #3179
@gmin7 made their first contribution in #3214

Full Changelog: v0.31.0...v0.31.1

Original source

Highlights

Initial version of QMMs for CUDA (#3160)

JACCL mesh bandwidth improvements (#3174)

Massive speedups for 3D convs (#3147)

Continued improvements to qqmm (#3106, #3022)

What's Changed

Patch bump by @angeloskath in #3102

is_available() should check the device index too by @andresy in #3107

Fix residency set with user provided buffer by @awni in #3108

Cleanup test_fast_sdpa.py by @zcbenz in #3112

[CUDA] Set current device before allocating memory by @zcbenz in #3110

Quantize module to QQLinear by @nastya236 in #3106

[CUDA] Use cuDNN SDPA for decoding when using fixed-size KV cache by @zcbenz in #3113

register pressure by @nastya236 in #3116

Fix precision in Metal fused attention by @awni in #3119

[CUDA] Attention sinks in cuDNN SDPA by @zcbenz in #3118

Fix donation in sdpa vector by @angeloskath in #3121

Manage stream placement in import function by @awni in #3127

fix: propagate quantization mode in QuantizedAllToShardedLinear / QuantizedShardedToAllLinear by @vskiwi in #3133

[featuring] - add hanning window function by @Vlor999 in #3124

feat: adding the hamming function by @Vlor999 in #3135

Tensor scale nvfp4 by @nastya236 in #3022

Fix fence synchronization accross command buffers by @awni in #3144

Export: preserve Dtype state values in export callback arguments by @skryl in #3145

[Metal] Fix 32-bit integer overflow in conv3d unfold kernel by @kellen-sun in #3143

[Metal][Performance] Add implicit matmul pathway for mx.conv3d by @belkakari in #3147

[Metal] Fix event leak by @awni in #3159

[CUDA] FPxINT quantized matmul for Hopper by @zcbenz in #3160

feat: implement mlx.core.blackman by @Vlor999 in #3136

Enable setting thread block cluster for Hopper and later by @zcbenz in #3168

[CUDA][NCCL] group split by @nastya236 in #3172

JACCL refactor and small update by @angeloskath in #3174

[CUDA] Heuristics for Hopper QMM by @zcbenz in #3173

Fix compile_fuse broadcast split aliasing bug by @robert-johansson in #3166

Enable passing in a GPU architecture string via env var by @angeloskath in #3176

Bump the minor version by @angeloskath in #3183

New Contributors

@vskiwi made their first contribution in #3133

@Vlor999 made their first contribution in #3124

@skryl made their first contribution in #3145

@kellen-sun made their first contribution in #3143

@belkakari made their first contribution in #3147

@robert-johansson made their first contribution in #3166

Full Changelog: v0.30.6...v0.31.0

Original source

Highlights

Much faster bandwidth with JACCL on macOS >= 26.3 (some numbers)

What's Changed

• patch by @awni in #3093
• Disable managed memory on WSL when concurrentManagedAccess is not supported by @jessegross in #3095
• Fix non simd f16 build by @awni in #3097
• Fix 2pass sdpa on < M2 by @awni in #3099
• JACCL update by @angeloskath in #3094
• Fix qmv_impl for small N by @manuelcandales in #3096
• Patch for multi device CUDA by @awni in #3100

New Contributors

• @manuelcandales made their first contribution in #3096

Full Changelog: v0.30.5...v0.30.6

Original source

What's Changed

• patch by @awni in #3074
• [CUDA] Fallback Event impl when there is no hardware cpu/gpu coherency by @zcbenz in #3070
• Tune CUDA gaph sizes on B200 and H100 by @awni in #3077
• [Docs] Simple example of using MLX distributed by @stefpi in #2973
• Use lower-right causal mask alignment consistently by @Anri-Lombard in #2967
• Fix ALiBi slopes for non-power-of-2 num_heads by @vovw in #3071
• More useful error for large indices by @awni in #3079
• Fix nax condition for iphone by @awni in #3083
• Fallback to pinned host memory when managed memory is not supported by @zcbenz in #3075
• Fix failing python tests on Windows by @zcbenz in #3076
• [Metal] Tune splitk gemm dispatch conditions and partition sizes by @awni in #3087
• Fix for NAX overflow. by @awni in #3092

New Contributors

• @stefpi made their first contribution in #2973
• @vovw made their first contribution in #3071

Full Changelog: v0.30.4...v0.30.5

Original source

Highlights

Metal: Much faster vector fused grouped-query attention for long context

CUDA: Several improvements to speed up LLM inference for CUDA backend

CUDA: Support for dense MoEs

CUDA: Better support for consumer GPUs (4090, 5090, RTX 6000, ...)

What's Changed

• patch bump for next release by @awni in #2991
• Fix fence by @awni in #2998
• Reverts changing the MLX_IBV_DEVICES to MLX_JACCL_DEVICES by @angeloskath in #2999
• fix distributed all_to_sharded bias shard axis from -2 to -1 by @gufengc in #2987
• Fix sharding of quantized models with non-power-of-2 bits by @kernelpool in #3006
• Update CCCL to v3.1.3 by @zcbenz in #3012
• Fix python package install path in stubgen by @zcbenz in #3009
• Type Enhancement for Func Transforms and Bug Fix by @XXXXRT666 in #3003
• Do not clear disk space in setup-linux by @zcbenz in #3013
• Do not give workflow boolean inputs default values by @zcbenz in #3014
• Fix negative dim indexing by @MillaFleurs in #2994
• Windows CI by @zcbenz in #3021
• Optimize erf function with expm1f in Metal backend by @bjornefisk in #3025
• [CUDA] Faster grouped mm by @zcbenz in #3011
• PR 3007 Fix Seg Fault by @MillaFleurs in #3008
• Use higher precision for linspace with double by @awni in #3029
• Handle data smaller than BUFFER_SIZE in jaccl recv by @rltakashige in #3033
• build 26.0 release in actions by @awni in #3035
• Remove xmlrunner from macOS CI by @zcbenz in #3032
• Columnwise quantize by @nastya236 in #2989
• Turn nccl_stub into a normal target by @zcbenz in #3037
• Use cuda::std for math ops by @zcbenz in #3041
• win: symbol exports and minor fixes by @dhiltgen in #3024
• CUDA gather mv by @angeloskath in #3039
• Link with prebuilt OpenBLAS and fix shared libs build on Windows by @zcbenz in #3036
• Allow take on empty array when it makes sense by @awni in #3046
• Add missing include to buffer_cache.h by @Anri-Lombard in #3053
• Build and test python package on Windows CI by @zcbenz in #3049
• Fix some MSVC compilation errors by @zcbenz in #3048
• Use C++20 by @zcbenz in #3050
• Faster two pass sdpa by @awni in #3023
• Find system-installed cuDNN on Windows by @zcbenz in #3052
• Fix some NVCC warnings when building CUDA backend with MSVC by @zcbenz in #3038
• Hide symbols by default for mac/linux by @zcbenz in #3057
• [CUDA] Fast sorting by @awni in #3060
• Fix flaky macOS test by @awni in #3063
• Update pre-commit hooks and versions for clang-format, black, and isort by @NripeshN in #3059
• GPU discovery by @dhiltgen in #3055
• Add NAX Split-K GEMM for large-K matmuls to improve performance by @hxu296 in #3018
• Improve CPU discovery by @dhiltgen in #3068
• Fix long cache file path on Windows by @zcbenz in #3065
• Better support consumer CUDA GPUs by @jessegross in #3056
• Delay load CUDA libs and resolve DLL paths at runtime by @zcbenz in #3061
• Do not require ConcurrentManagedAccess when not used by @zcbenz in #3062
• Fp qmv by @awni in #2984
• remove thrust by @awni in #3067

New Contributors

• @gufengc made their first contribution in #2987
• @kernelpool made their first contribution in #3006
• @bjornefisk made their first contribution in #3025
• @rltakashige made their first contribution in #3033
• @dhiltgen made their first contribution in #3024
• @hxu296 made their first contribution in #3018
• @jessegross made their first contribution in #3056

Full Changelog: v0.30.3...v0.30.4

Original source

Highlights

Support nvfp4 and mxfp8 quantized ops on Metal

Support nvfp4 and mxfp8 quantized-quantized matrix-matrix multiplication on CUDA

What's Changed

Bump the patch version by @angeloskath in #2922

Faster copy for col contig to row contig by @awni in #2917

Fix cuda release by @awni in #2925

Metal logging by @CC-Yeh in #2904

fix cuda release part 2 by @awni in #2926

new[CI]: add linux sanitizer tests by @incertum in #2860

patch bump by @awni in #2927

Fix CUDA pypi release by @awni in #2929

Move allocate_workspace to cuda/utils.h by @zcbenz in #2923

Allow dry run for PyPI release workflow by @zcbenz in #2928

Set rpath with cmake for CUDA build by @zcbenz in #2932

Fix nightly build by @zcbenz in #2933

Set install rpath of python bindings with cmake by @zcbenz in #2934

Fix pid in local launch by @angeloskath in #2936

Make CUDA CI run faster by @zcbenz in #2939

refactor: use perf_counter for accurate benchmarking by @Satyam12singh in #2940

Fix for non row-contig scales by @awni in #2941

Fix stubgen by @zcbenz in #2942

ci: add macOS 26 target by @madrob in #2937

Fix float64 size in data_types.rst by @pdevine in #2948

Fixes in mlx.distributed_config by @angeloskath in #2947

Metal/CPU nvfp4 and mxfp8 by @awni in #2946

[CUDA] Implement gather_mm_rhs by @zcbenz in #2902

Fetch nanobind with cmake by @zcbenz in #2949

refactor: use time.perf_counter for consistent and accurate benchmarking by @Satyam12singh in #2943

BUG FIX - Addition of missing parameter in random::uniform by @hwiesmann in #2963

Fix doc issues in mlx.nn.init.he_normal and mlx.nn.hard_tanh by @Redempt1onzzZZ in #2968

fix numpy dtype bug by @awni in #2960

QQ linear by @nastya236 in #2931

fix array allocator with user buffer and deleter by @andresy in #2971

Swizzle scales by @nastya236 in #2979

Fix grid_dim_x calculations by @CC-Yeh in #2980

Add asarray to array_namespace by @Anri-Lombard in #2966

fix doc by @CC-Yeh in #2988

replace MLX_IBV_COORDINATOR with MLX_JACCL_COORDINATOR by @Evanev7 in #2986

Fix RandomBits::is_equivalent to include width by @MillaFleurs in #2978

Don't try to use NAX at run-time if kernels aren't there by @awni in #2982

Expose to/from fp8 in Python and don't auto-convert fp8 when loading from safetensors by @awni in #2985

Allow some non 2D inputs in qqmm by @awni in #2981

New Contributors

@pdevine made their first contribution in #2948

@hwiesmann made their first contribution in #2963

@Anri-Lombard made their first contribution in #2966

@Evanev7 made their first contribution in #2986

@MillaFleurs made their first contribution in #2978

Full Changelog: v0.30.1...v0.30.3

Original source