only interrupt during an eval

.circleci/config.yml

@@ -24,8 +24,8 @@ jobs:
         type: boolean
         default: false
     macos:
-      xcode: "16.2.0"
-    resource_class: m2pro.medium
+      xcode: "15.2.0"
+    resource_class: macos.m1.medium.gen1
     steps:
       - checkout
       - run:
@@ -89,14 +89,15 @@ jobs:
             pip install numpy
             sudo apt-get update
             sudo apt-get install libblas-dev liblapack-dev liblapacke-dev
-            sudo apt-get install openmpi-bin openmpi-common libopenmpi-dev
       - run:
           name: Install Python package
           command: |
-            CMAKE_ARGS="-DMLX_BUILD_METAL=OFF" \
+            CMAKE_ARGS="-DMLX_BUILD_METAL=OFF
+              CMAKE_COMPILE_WARNING_AS_ERROR=ON" \
               CMAKE_BUILD_PARALLEL_LEVEL=`nproc` \
               python3 setup.py build_ext --inplace
-            CMAKE_ARGS="-DMLX_BUILD_METAL=OFF" \
+            CMAKE_ARGS="-DMLX_BUILD_METAL=OFF \
+              CMAKE_COMPILE_WARNING_AS_ERROR=ON" \
               CMAKE_BUILD_PARALLEL_LEVEL=`nproc` \
               python3 setup.py develop
       - run:
@@ -109,8 +110,6 @@ jobs:
           name: Run Python tests
           command: |
             python3 -m unittest discover python/tests -v
-            mpirun --bind-to none -host localhost:8 -np 8 python python/tests/mpi_test_distributed.py
-            mlx.launch --verbose -n 8 python/tests/ring_test_distributed.py
       - run:
           name: Build CPP only
           command: |
@@ -125,15 +124,10 @@ jobs:
     parameters:
       xcode_version:
         type: string
-        default: "16.2.0"
-      macosx_deployment_target:
-        type: string
-        default: ""
+        default: "15.2.0"
     macos:
       xcode: << parameters.xcode_version >>
-    environment:
-      MACOSX_DEPLOYMENT_TARGET: << parameters.macosx_deployment_target >>
-    resource_class: m2pro.medium
+    resource_class: macos.m1.medium.gen1
     steps:
       - checkout
       - run:
@@ -155,7 +149,7 @@ jobs:
           command: |
             source env/bin/activate
             DEBUG=1 CMAKE_BUILD_PARALLEL_LEVEL=`sysctl -n hw.ncpu` \
-            CMAKE_ARGS="-DCMAKE_COMPILE_WARNING_AS_ERROR=ON" \
+            CMAKE_ARGS="CMAKE_COMPILE_WARNING_AS_ERROR=ON" \
               pip install -e . -v
       - run:
           name: Generate package stubs
@@ -219,18 +213,13 @@ jobs:
         default: "3.9"
       xcode_version:
         type: string
-        default: "16.2.0"
+        default: "15.2.0"
       build_env:
         type: string
         default: ""
-      macosx_deployment_target:
-        type: string
-        default: ""
     macos:
       xcode: << parameters.xcode_version >>
-    resource_class: m2pro.medium
-    environment:
-      MACOSX_DEPLOYMENT_TARGET: << parameters.macosx_deployment_target >>
+    resource_class: macos.m1.medium.gen1
     steps:
       - checkout
       - run:
@@ -251,7 +240,7 @@ jobs:
           name: Install Python package
           command: |
             source env/bin/activate
-            env -u MACOSX_DEPLOYMENT_TARGET DEV_RELEASE=1 \
+            DEV_RELEASE=1 \
               CMAKE_BUILD_PARALLEL_LEVEL=`sysctl -n hw.ncpu` \
               pip install . -v
       - run:
@@ -346,7 +335,7 @@ workflows:
       - mac_build_and_test:
           matrix:
             parameters:
-              macosx_deployment_target: ["13.5", "14.0"]
+              xcode_version: ["15.0.0", "15.2.0", "16.0.0"]
       - linux_build_and_test
       - build_documentation 
 
@@ -366,70 +355,8 @@ workflows:
           matrix:
             parameters:
               python_version: ["3.9", "3.10", "3.11", "3.12", "3.13"]
-              macosx_deployment_target: ["13.5", "14.0", "15.0"]
+              xcode_version: ["15.0.0", "15.2.0"]
               build_env: ["PYPI_RELEASE=1"]
-              xcode_version: ["16.2.0", "15.0.0"]
-            exclude:
-              - macosx_deployment_target: "13.5"
-                xcode_version: "16.2.0"
-                python_version: "3.9"
-                build_env: "PYPI_RELEASE=1"
-              - macosx_deployment_target: "13.5"
-                xcode_version: "16.2.0"
-                python_version: "3.10"
-                build_env: "PYPI_RELEASE=1"
-              - macosx_deployment_target: "13.5"
-                xcode_version: "16.2.0"
-                python_version: "3.11"
-                build_env: "PYPI_RELEASE=1"
-              - macosx_deployment_target: "13.5"
-                xcode_version: "16.2.0"
-                python_version: "3.12"
-                build_env: "PYPI_RELEASE=1"
-              - macosx_deployment_target: "13.5"
-                xcode_version: "16.2.0"
-                python_version: "3.13"
-                build_env: "PYPI_RELEASE=1"
-              - macosx_deployment_target: "14.0"
-                xcode_version: "15.0.0"
-                python_version: "3.9"
-                build_env: "PYPI_RELEASE=1"
-              - macosx_deployment_target: "14.0"
-                xcode_version: "15.0.0"
-                python_version: "3.10"
-                build_env: "PYPI_RELEASE=1"
-              - macosx_deployment_target: "14.0"
-                xcode_version: "15.0.0"
-                python_version: "3.11"
-                build_env: "PYPI_RELEASE=1"
-              - macosx_deployment_target: "14.0"
-                xcode_version: "15.0.0"
-                python_version: "3.12"
-                build_env: "PYPI_RELEASE=1"
-              - macosx_deployment_target: "14.0"
-                xcode_version: "15.0.0"
-                python_version: "3.13"
-                build_env: "PYPI_RELEASE=1"
-              - macosx_deployment_target: "15.0"
-                xcode_version: "15.0.0"
-                python_version: "3.9"
-                build_env: "PYPI_RELEASE=1"
-              - macosx_deployment_target: "15.0"
-                xcode_version: "15.0.0"
-                python_version: "3.10"
-                build_env: "PYPI_RELEASE=1"
-              - macosx_deployment_target: "15.0"
-                xcode_version: "15.0.0"
-                python_version: "3.11"
-                build_env: "PYPI_RELEASE=1"
-              - macosx_deployment_target: "15.0"
-                xcode_version: "15.0.0"
-                python_version: "3.12"
-                build_env: "PYPI_RELEASE=1"
-              - macosx_deployment_target: "15.0"
-                xcode_version: "15.0.0"
-                python_version: "3.13"
-                build_env: "PYPI_RELEASE=1"
       - build_documentation:
           filters:
             tags:
@@ -452,7 +379,7 @@ workflows:
           requires: [ hold ]
           matrix:
             parameters:
-              macosx_deployment_target: ["13.5", "14.0"]
+              xcode_version: ["15.0.0", "15.2.0", "16.0.0"]
       - linux_build_and_test:
           requires: [ hold ]
   nightly_build:
@@ -465,54 +392,7 @@ workflows:
           matrix:
             parameters:
               python_version: ["3.9", "3.10", "3.11", "3.12", "3.13"]
-              macosx_deployment_target: ["13.5", "14.0", "15.0"]
-              xcode_version: ["16.2.0", "15.0.0"]
-            exclude:
-              - macosx_deployment_target: "13.5"
-                xcode_version: "16.2.0"
-                python_version: "3.9"
-              - macosx_deployment_target: "13.5"
-                xcode_version: "16.2.0"
-                python_version: "3.10"
-              - macosx_deployment_target: "13.5"
-                xcode_version: "16.2.0"
-                python_version: "3.11"
-              - macosx_deployment_target: "13.5"
-                xcode_version: "16.2.0"
-                python_version: "3.12"
-              - macosx_deployment_target: "13.5"
-                xcode_version: "16.2.0"
-                python_version: "3.13"
-              - macosx_deployment_target: "14.0"
-                xcode_version: "15.0.0"
-                python_version: "3.9"
-              - macosx_deployment_target: "14.0"
-                xcode_version: "15.0.0"
-                python_version: "3.10"
-              - macosx_deployment_target: "14.0"
-                xcode_version: "15.0.0"
-                python_version: "3.11"
-              - macosx_deployment_target: "14.0"
-                xcode_version: "15.0.0"
-                python_version: "3.12"
-              - macosx_deployment_target: "14.0"
-                xcode_version: "15.0.0"
-                python_version: "3.13"
-              - macosx_deployment_target: "15.0"
-                xcode_version: "15.0.0"
-                python_version: "3.9"
-              - macosx_deployment_target: "15.0"
-                xcode_version: "15.0.0"
-                python_version: "3.10"
-              - macosx_deployment_target: "15.0"
-                xcode_version: "15.0.0"
-                python_version: "3.11"
-              - macosx_deployment_target: "15.0"
-                xcode_version: "15.0.0"
-                python_version: "3.12"
-              - macosx_deployment_target: "15.0"
-                xcode_version: "15.0.0"
-                python_version: "3.13"
+              xcode_version: ["15.0.0", "15.2.0"]
   weekly_build:
     when:
       and:
@@ -523,70 +403,8 @@ workflows:
           matrix:
             parameters:
               python_version: ["3.9", "3.10", "3.11", "3.12", "3.13"]
-              macosx_deployment_target: ["13.5", "14.0", "15.0"]
+              xcode_version: ["15.0.0", "15.2.0", "16.0.0"]
               build_env: ["DEV_RELEASE=1"]
-              xcode_version: ["16.2.0", "15.0.0"]
-            exclude:
-              - macosx_deployment_target: "13.5"
-                xcode_version: "16.2.0"
-                python_version: "3.9"
-                build_env: "DEV_RELEASE=1"
-              - macosx_deployment_target: "13.5"
-                xcode_version: "16.2.0"
-                python_version: "3.10"
-                build_env: "DEV_RELEASE=1"
-              - macosx_deployment_target: "13.5"
-                xcode_version: "16.2.0"
-                python_version: "3.11"
-                build_env: "DEV_RELEASE=1"
-              - macosx_deployment_target: "13.5"
-                xcode_version: "16.2.0"
-                python_version: "3.12"
-                build_env: "DEV_RELEASE=1"
-              - macosx_deployment_target: "13.5"
-                xcode_version: "16.2.0"
-                python_version: "3.13"
-                build_env: "DEV_RELEASE=1"
-              - macosx_deployment_target: "14.0"
-                xcode_version: "15.0.0"
-                python_version: "3.9"
-                build_env: "DEV_RELEASE=1"
-              - macosx_deployment_target: "14.0"
-                xcode_version: "15.0.0"
-                python_version: "3.10"
-                build_env: "DEV_RELEASE=1"
-              - macosx_deployment_target: "14.0"
-                xcode_version: "15.0.0"
-                python_version: "3.11"
-                build_env: "DEV_RELEASE=1"
-              - macosx_deployment_target: "14.0"
-                xcode_version: "15.0.0"
-                python_version: "3.12"
-                build_env: "DEV_RELEASE=1"
-              - macosx_deployment_target: "14.0"
-                xcode_version: "15.0.0"
-                python_version: "3.13"
-                build_env: "DEV_RELEASE=1"
-              - macosx_deployment_target: "15.0"
-                xcode_version: "15.0.0"
-                python_version: "3.9"
-                build_env: "DEV_RELEASE=1"
-              - macosx_deployment_target: "15.0"
-                xcode_version: "15.0.0"
-                python_version: "3.10"
-                build_env: "DEV_RELEASE=1"
-              - macosx_deployment_target: "15.0"
-                xcode_version: "15.0.0"
-                python_version: "3.11"
-                build_env: "DEV_RELEASE=1"
-              - macosx_deployment_target: "15.0"
-                xcode_version: "15.0.0"
-                python_version: "3.12"
-                build_env: "DEV_RELEASE=1"
-              - macosx_deployment_target: "15.0"
-                xcode_version: "15.0.0"
-                python_version: "3.13"
-                build_env: "DEV_RELEASE=1"
   linux_test_release:
     when:
       and:

CMakeLists.txt

@@ -212,6 +212,24 @@ else()
   set(MLX_BUILD_ACCELERATE OFF)
 endif()
 
+find_package(MPI)
+if(MPI_FOUND)
+  execute_process(
+    COMMAND zsh "-c" "mpirun --version"
+    OUTPUT_VARIABLE MPI_VERSION
+    ERROR_QUIET)
+  if(${MPI_VERSION} MATCHES ".*Open MPI.*")
+    target_include_directories(mlx PRIVATE ${MPI_INCLUDE_PATH})
+  elseif(MPI_VERSION STREQUAL "")
+    set(MPI_FOUND FALSE)
+    message(
+      WARNING "MPI found but mpirun is not available. Building without MPI.")
+  else()
+    set(MPI_FOUND FALSE)
+    message(WARNING "MPI which is not OpenMPI found. Building without MPI.")
+  endif()
+endif()
+
 message(STATUS "Downloading json")
 FetchContent_Declare(
   json

CONTRIBUTING.md

@@ -5,26 +5,26 @@ possible.
 
 ## Pull Requests
 
-1. Fork and submit pull requests to the repo.
+1. Fork and submit pull requests to the repo. 
 2. If you've added code that should be tested, add tests.
 3. If a change is likely to impact efficiency, run some of the benchmarks before
    and after the change. Examples of benchmarks can be found in `benchmarks/python/`.
 4. If you've changed APIs, update the documentation.
-5. Every PR should have passing tests and at least one review.
+5. Every PR should have passing tests and at least one review. 
 6. For code formatting install `pre-commit` using something like `pip install pre-commit` and run `pre-commit install`.
    This should install hooks for running `black` and `clang-format` to ensure
    consistent style for C++ and python code.
-
+ 
    You can also run the formatters manually as follows:
-
-   ```shell
-   clang-format -i file.cpp
-   ```
-
-   ```shell
-   black file.py
-   ```
-
+ 
+     ```
+     clang-format -i file.cpp
+     ```
+ 
+     ```
+     black file.py
+     ```
+ 
    or run `pre-commit run --all-files` to check all files in the repo.
 
 ## Issues

benchmarks/python/blas/bench_gemm.py

@@ -157,7 +157,7 @@ def bench_shape(B, M, N, K, np_dtype, transpose="nn"):
 
 
 def get_gflop_count(B, M, N, K):
-    return float(2.0 * N_iter_bench * N_iter_func * B * M * N * K) / float(1000.0**3)
+    return float(2.0 * N_iter_bench * N_iter_func * B * M * N * K) / float(1024.0**3)
 
 
 if __name__ == "__main__":
@@ -175,8 +175,6 @@ if __name__ == "__main__":
         (1, 4096, 4096, 4096),
     )
 
-    print(f"  B,    M,    N,     K,   dtype,  t, gflops_pt, gflops_mx, diff%")
-
     for dtype in dtypes:
         for transpose in transposes:
             for B, M, N, K in shapes:
@@ -189,7 +187,7 @@ if __name__ == "__main__":
                 diff = gflops_mx / gflops_pt - 1.0
 
                 print(
-                    f"{B:3d}, {M:4d}, {N:4d}, {K:5d}, {dtype}, {transpose},  {gflops_pt:8.2f},  {gflops_mx:8.2f}, {100. * diff:+5.2f}%"
+                    f"{B:3d}, {M:4d}, {N:4d}, {K:4d}, {dtype}, {transpose}, {gflops_pt:05.3f}, {gflops_mx:05.3f}, {100. * diff:+5.2f}%"
                 )
                 if gflops_pt >= 2.0 * gflops_mx:
                     print("ATTENTION ^^^^^^^")

benchmarks/python/gather_mm_bench.py

@@ -1,74 +0,0 @@
-# Copyright © 2023-2024 Apple Inc.
-
-import mlx.core as mx
-from time_utils import time_fn
-
-N = 1024
-D = 1024
-M = 1024
-E = 32
-I = 4
-
-
-def gather_sort(x, indices):
-    N, M = indices.shape
-    indices = indices.flatten()
-    order = mx.argsort(indices)
-    inv_order = mx.argsort(order)
-    return x.flatten(0, -3)[order // M], indices[order], inv_order
-
-
-def scatter_unsort(x, inv_order, shape=None):
-    x = x[inv_order]
-    if shape is not None:
-        x = mx.unflatten(x, 0, shape)
-    return x
-
-
-def gather_mm_simulate(x, w, indices):
-    x, idx, inv_order = gather_sort(x, indices)
-    for i in range(2):
-        y = mx.concatenate([x[i] @ w[j].T for i, j in enumerate(idx.tolist())], axis=0)
-        x = y[:, None]
-    x = scatter_unsort(x, inv_order, indices.shape)
-    return x
-
-
-def time_gather_mm():
-    x = mx.random.normal((N, 1, 1, D)) / 1024**0.5
-    w1 = mx.random.normal((E, M, D)) / 1024**0.5
-    w2 = mx.random.normal((E, D, M)) / 1024**0.5
-    indices = (mx.random.uniform(shape=(N, I)) * E).astype(mx.uint32)
-    sorted_indices = mx.sort(indices.flatten()).reshape(N, I)
-    mx.eval(x, w1, w2, indices, sorted_indices)
-
-    def gather_mm(x, w1, w2, indices, sort):
-        idx = indices
-        inv_order = None
-        if sort:
-            x, idx, inv_order = gather_sort(x, indices)
-        x = mx.gather_mm(x, w1.swapaxes(-1, -2), rhs_indices=idx, sorted_indices=sort)
-        x = mx.gather_mm(x, w2.swapaxes(-1, -2), rhs_indices=idx, sorted_indices=sort)
-        if sort:
-            x = scatter_unsort(x, inv_order, indices.shape)
-        return x
-
-    time_fn(gather_mm, x, w1, w2, indices, False)
-    time_fn(gather_mm, x, w1, w2, sorted_indices, False)
-    time_fn(gather_mm, x, w1, w2, indices, True)
-
-    x = mx.random.normal((N * I, D)) / 1024**0.5
-    w1 = mx.random.normal((M, D)) / 1024**0.5
-    w2 = mx.random.normal((D, M)) / 1024**0.5
-    mx.eval(x, w1, w2)
-
-    def equivalent_matmul(x, w1, w2):
-        x = x @ w1.T
-        x = x @ w2.T
-        return x
-
-    time_fn(equivalent_matmul, x, w1, w2)
-
-
-if __name__ == "__main__":
-    time_gather_mm()

benchmarks/python/sdpa_bench.py

@@ -28,34 +28,11 @@ def bench(f, *args):
     return (e - s) * 1e-9
 
 
-def prepare_inputs(B, qL, kL, D, qH, kH, mask, transpose, dtype):
-    np_dtype = getattr(np, dtype)
-
-    shape_q = (B, qL, qH, D) if transpose else (B, qH, qL, D)
-    shape_kv = (B, kL, kH, D) if transpose else (B, kH, kL, D)
-
-    scale = 1.0 / math.sqrt(D)
-
-    q_np = np.random.normal(0.0, 1.0, shape_q).astype(np_dtype)
-    k_np = np.random.normal(0.0, scale, shape_kv).astype(np_dtype)
-    v_np = np.random.normal(0.0, scale, shape_kv).astype(np_dtype)
-
-    q_mx = mx.array(q_np)
-    k_mx = mx.array(k_np)
-    v_mx = mx.array(v_np)
-
-    if mask is not None:
-        if mask == "additive":
-            mask_np = np.random.normal(0.0, 1.0, (B, qH, qL, kL)).astype(np_dtype)
-            mask = mx.array(mask_np)
-        elif mask == "bool":
-            mask_np = np.random.uniform(0.0, 1.0, (B, qH, qL, kL)) < 0.5
-            mask = mx.array(mask_np)
-
-    return q_mx, k_mx, v_mx, scale, mask
+def mlx_sdpa_fused_inner(q, k, v, scale):
+    return mx.fast.scaled_dot_product_attention(q, k, v, scale=scale, mask=None)
 
 
-def mlx_ref_attn(q, k, v, scale=1.0, mask=None):
+def mlx_sdpa_unfused_inner(q, k, v, scale, f32softmax=False):
     q_dtype = q.dtype
     q = q * mx.array(scale, q_dtype)
     n_q_heads = q.shape[-3]
@@ -64,7 +41,6 @@ def mlx_ref_attn(q, k, v, scale=1.0, mask=None):
 
     B = q.shape[0]
     L = q.shape[2]
-    kL = k.shape[2]
 
     if n_repeats > 1:
         q = mx.reshape(q, [B, n_kv_heads, n_repeats, L, -1])
@@ -72,27 +48,10 @@ def mlx_ref_attn(q, k, v, scale=1.0, mask=None):
         v = mx.expand_dims(v, 2)
 
     scores = q @ mx.swapaxes(k, -1, -2)
-
-    if mask is not None:
-
-        if mask == "causal":
-            q_offset = max(0, kL - L)
-            q_indices = mx.arange(q_offset, q_offset + L)
-            k_indices = mx.arange(kL)
-            mask = q_indices[:, None] >= k_indices[None]
-
-        if n_repeats > 1 and mask.ndim >= 3:
-            if mask.shape[-3] == 1:
-                mask = mx.expand_dims(mask, -3)
-            else:
-                mask = mx.unflatten(mask, -3, (n_kv_heads, n_repeats))
-
-        if mask.dtype == mx.bool_:
-            scores = mx.where(mask, scores, -np.float32(np.inf))
-        else:
-            scores += mask
-
-    scores = mx.softmax(scores, axis=-1, precise=True)
+    if f32softmax:
+        scores = mx.softmax(scores.astype(mx.float32), axis=-1).astype(q_dtype)
+    else:
+        scores = mx.softmax(scores, axis=-1)
 
     out = scores @ v
     if n_repeats > 1:
@@ -101,55 +60,74 @@ def mlx_ref_attn(q, k, v, scale=1.0, mask=None):
     return out
 
 
-def mlx_fused_attn(q, k, v, scale, mask):
-    return mx.fast.scaled_dot_product_attention(q, k, v, scale=scale, mask=mask)
-
-
-def do_attention(f, q, k, v, scale, mask=None, transpose=False):
-    if transpose:
-        q_t = mx.transpose(q, (0, 2, 1, 3))
-        k_t = mx.transpose(k, (0, 2, 1, 3))
-        v_t = mx.transpose(v, (0, 2, 1, 3))
-        o_t = f(q_t, k_t, v_t, scale=scale, mask=mask)
-        return mx.transpose(o_t, (0, 2, 1, 3))
-    else:
-        return f(q, k, v, scale=scale, mask=mask)
-
-
-def do_attention_bench(f, q, k, v, scale, mask=None, transpose=False):
+def mlx_spda_unfused(q, k, v, scale, transpose):
     q_out = q
+    if transpose:
+        k = mx.transpose(k, (0, 2, 1, 3))
+        v = mx.transpose(v, (0, 2, 1, 3))
 
     for i in range(N_iter_func):
-        q_out = do_attention(f, q_out, k, v, scale, mask=mask, transpose=transpose)
+        if transpose:
+            q_out = mx.transpose(q_out, (0, 2, 1, 3))
+        q_out = mlx_sdpa_unfused_inner(q_out, k, v, scale)
+        if transpose:
+            q_out = mx.transpose(q_out, (0, 2, 1, 3))
 
     mx.eval(q_out)
     return q_out
 
 
-def bench_shape(
-    B, qsl, ksl, head_dim, n_q_heads, n_kv_heads, dtype, transpose=True, mask_in=None
-):
-    q_mx, k_mx, v_mx, scale, mask = prepare_inputs(
-        B, qsl, ksl, head_dim, n_q_heads, n_kv_heads, mask_in, transpose, dtype
+def mlx_spda_fused(q, k, v, scale, transpose):
+    q_out = q
+    if transpose:
+        k = mx.transpose(k, (0, 2, 1, 3))
+        v = mx.transpose(v, (0, 2, 1, 3))
+
+    for i in range(N_iter_func):
+        if transpose:
+            q_out = mx.transpose(q_out, (0, 2, 1, 3))
+        q_out = mlx_sdpa_fused_inner(q_out, k, v, scale)
+        if transpose:
+            q_out = mx.transpose(q_out, (0, 2, 1, 3))
+
+    mx.eval(q_out)
+    return q_out
+
+
+def bench_shape(B, qsl, ksl, head_dim, n_q_heads, n_kv_heads, np_dtype, transpose=True):
+    shape_q = (
+        (B, qsl, n_q_heads, head_dim) if transpose else (B, n_q_heads, qsl, head_dim)
+    )
+    shape_kv = (
+        (B, ksl, n_kv_heads, head_dim) if transpose else (B, n_kv_heads, ksl, head_dim)
     )
 
-    time_mlx_unfused = bench(
-        do_attention_bench, mlx_ref_attn, q_mx, k_mx, v_mx, scale, mask, transpose
-    )
-    time_mlx_fused = bench(
-        do_attention_bench, mlx_fused_attn, q_mx, k_mx, v_mx, scale, mask, transpose
-    )
+    q_np = np.random.normal(0.0, 1.0 / math.sqrt(head_dim), shape_q).astype(np_dtype)
+    k_np = np.random.normal(0.0, 1.0 / math.sqrt(head_dim), shape_kv).astype(np_dtype)
+    v_np = np.random.normal(0.0, 1.0 / math.sqrt(head_dim), shape_kv).astype(np_dtype)
 
-    o_mlx_fused = do_attention(mlx_ref_attn, q_mx, k_mx, v_mx, scale, mask, transpose)
-    o_mlx_unfused = do_attention(
-        mlx_fused_attn, q_mx, k_mx, v_mx, scale, mask, transpose
-    )
+    scale = math.sqrt(1.0 / head_dim)
 
-    atol = 1e-5 if dtype == "float32" else 2e-4
+    q_mx = mx.array(q_np)
+    k_mx = mx.array(k_np)
+    v_mx = mx.array(v_np)
 
-    if not mx.allclose(o_mlx_fused, o_mlx_unfused, atol=atol, rtol=atol):
+    time_mlx_unfused = bench(mlx_spda_unfused, q_mx, k_mx, v_mx, scale, transpose)
+    time_mlx_fused = bench(mlx_spda_fused, q_mx, k_mx, v_mx, scale, transpose)
+
+    if transpose:
+        q_mx = mx.transpose(q_mx, (0, 2, 1, 3))
+        k_mx = mx.transpose(k_mx, (0, 2, 1, 3))
+        v_mx = mx.transpose(v_mx, (0, 2, 1, 3))
+
+    o_mlx_fused = mlx_sdpa_fused_inner(q_mx, k_mx, v_mx, scale)
+    o_mlx_unfused = mlx_sdpa_unfused_inner(q_mx, k_mx, v_mx, scale, f32softmax=True)
+
+    atol = 1e-5 if np_dtype == np.float32 else 1e-4
+
+    if not mx.allclose(o_mlx_fused, o_mlx_unfused, atol=atol):
         print(
-            f"Failed at (B: {B}, qsl: {qsl}, ksl: {ksl}, head_dim: {head_dim}, n_qh: {n_q_heads}, n_kvh: {n_kv_heads}, mask: {mask_in}) [tpose = {transpose}] with max(|a - b|) = {mx.max(mx.abs(o_mlx_unfused - o_mlx_fused)):3.2e}"
+            f"Failed at (B: {B}, qsl: {qsl}, ksl: {ksl}, head_dim: {head_dim}, n_qh: {n_q_heads}, n_kvh: {n_kv_heads}) [tpose = {transpose}] with max(|a - b|) = {mx.max(mx.abs(o_mlx_unfused - o_mlx_fused)):3.2e}"
         )
 
     return time_mlx_fused, time_mlx_unfused
@@ -173,51 +151,39 @@ if __name__ == "__main__":
           (  1,   128,   128,       64,   32,    32),
           (  1,   256,   256,       64,   32,    32),
           (  1,   512,   512,       64,   32,    32),
-          (  1,  1024,  1024,       64,   32,     8),
-          (  1,  2048,  2048,       64,   32,     8),
-          (  1,  4096,  4096,       64,   32,     8),
+          (  1,  1024,  1024,       64,   32,    32),
+          (  1,  2048,  2048,       64,   32,    32),
+          (  1,  4096,  4096,       64,   32,    32),
     )
 
     shapes_80 = (
         # (  B,   qsl,   ksl, head_dim, n_qh, n_kvh)
-          (  1,  1024,  1024,       80,   32,     8),
-          (  1,  2048,  2048,       80,   32,     8),
-          (  1,  4096,  4096,       80,   32,     8),
+          (  1,  1024,  1024,       80,   32,    32),
+          (  1,  2048,  2048,       80,   32,    32),
+          (  1,  4096,  4096,       80,   32,    32),
     )
 
     shapes_128 = (
         # (  B,   qsl,   ksl, head_dim, n_qh, n_kvh)
-          (  1,  1024,  1024,      128,   32,     8),
-          (  1,  2048,  2048,      128,   32,     8),
-          (  1,  4096,  4096,      128,   32,     8),
+          (  1,  1024,  1024,      128,   32,    32),
+          (  1,  2048,  2048,      128,   32,    32),
+          (  1,  4096,  4096,      128,   32,    32),
     )
     # fmt: on
 
     shapes = shapes_64 + shapes_80 + shapes_128
 
-    masks = [None, "bool", "causal"]
-
-    print(
-        "  B,   qsl,   ksl, hdim, n_qh, n_kvh, t,   dtype,     mask, t_unfs, t_fuse, diff%"
-    )
+    print("  B,   qsl,   ksl, hdim, n_qh, n_kvh, tpose,   dtype, t_unfs, t_fuse, diff%")
 
     for dtype in dtypes:
         for transpose in transposes:
             for B, qsl, ksl, head_dim, n_q_heads, n_kv_heads in shapes:
-                for mask_in in masks:
-                    time_mlx_fused, time_mlx_unfused = bench_shape(
-                        B,
-                        qsl,
-                        ksl,
-                        head_dim,
-                        n_q_heads,
-                        n_kv_heads,
-                        dtype,
-                        transpose,
-                        mask_in,
-                    )
-                    diff = time_mlx_unfused / time_mlx_fused - 1.0
-                    t_str = 1 if transpose else 0
-                    print(
-                        f"{B:3d}, {qsl:5d}, {ksl:5d}, {head_dim:4d}, {n_q_heads:4d}, {n_kv_heads:5d}, {t_str:1d}, {dtype}, {str(mask_in):>8}, {time_mlx_unfused: 2.3f}, {time_mlx_fused: 2.3f}, {100. * diff:+5.2f}%"
-                    )
+                np_dtype = getattr(np, dtype)
+                time_mlx_fused, time_mlx_unfused = bench_shape(
+                    B, qsl, ksl, head_dim, n_q_heads, n_kv_heads, np_dtype, transpose
+                )
+                diff = time_mlx_unfused / time_mlx_fused - 1.0
+                t_str = 1 if transpose else 0
+                print(
+                    f"{B:3d}, {qsl:5d}, {ksl:5d}, {head_dim:4d}, {n_q_heads:4d}, {n_kv_heads:5d}, {t_str:5d}, {dtype}, {time_mlx_unfused: 2.3f}, {time_mlx_fused: 2.3f}, {100. * diff:+5.2f}%"
+                )

docs/Doxyfile

@@ -13,7 +13,7 @@ EXCLUDE_PATTERNS       = */private/*
 CREATE_SUBDIRS         = NO
 FULL_PATH_NAMES        = YES
 RECURSIVE              = YES
-GENERATE_HTML          = NO
+GENERATE_HTML          = YES
 GENERATE_LATEX         = NO
 GENERATE_XML           = YES
 XML_PROGRAMLISTING     = YES

docs/src/dev/extensions.rst

@@ -93,9 +93,9 @@ Primitives
 ^^^^^^^^^^^
 
 A :class:`Primitive` is part of the computation graph of an :class:`array`. It
-defines how to create output arrays given input arrays. Further, a
+defines how to create outputs arrays given a input arrays. Further, a
 :class:`Primitive` has methods to run on the CPU or GPU and for function
-transformations such as ``vjp`` and ``jvp``.  Let's go back to our example to be
+transformations such as ``vjp`` and ``jvp``.  Lets go back to our example to be
 more concrete:
 
 .. code-block:: C++
@@ -128,7 +128,7 @@ more concrete:
         /** The vector-Jacobian product. */
         std::vector<array> vjp(
             const std::vector<array>& primals,
-            const std::vector<array>& cotangents,
+            const array& cotan,
             const std::vector<int>& argnums,
             const std::vector<array>& outputs) override;
 
@@ -247,7 +247,9 @@ point-wise. This is captured in the templated function :meth:`axpby_impl`.
       float alpha_,
       float beta_,
       mx::Stream stream) {
-    out.set_data(mx::allocator::malloc(out.nbytes()));
+    // Allocate the output with `malloc_or_wait` which synchronously allocates
+    // memory, potentially waiting if the system is under memory pressure
+    out.set_data(mx::allocator::malloc_or_wait(out.nbytes()));
 
     // Get the CPU command encoder and register input and output arrays
     auto& encoder = mx::cpu::get_command_encoder(stream);
@@ -391,7 +393,7 @@ below.
         auto& d = metal::device(s.device);
 
         // Allocate output memory
-        out.set_data(allocator::malloc(out.nbytes()));
+        out.set_data(allocator::malloc_or_wait(out.nbytes()));
 
         // Resolve name of kernel
         std::ostringstream kname;
@@ -469,7 +471,7 @@ one we just defined:
             const std::vector<array>& tangents,
             const std::vector<int>& argnums) {
         // Forward mode diff that pushes along the tangents
-        // The jvp transform on the primitive can be built with ops
+        // The jvp transform on the primitive can built with ops
         // that are scheduled on the same stream as the primitive
 
         // If argnums = {0}, we only push along x in which case the
@@ -481,7 +483,7 @@ one we just defined:
             auto scale_arr = array(scale, tangents[0].dtype());
             return {multiply(scale_arr, tangents[0], stream())};
         }
-        // If argnums = {0, 1}, we take contributions from both
+        // If, argnums = {0, 1}, we take contributions from both
         // which gives us jvp = tangent_x * alpha + tangent_y * beta
         else {
             return {axpby(tangents[0], tangents[1], alpha_, beta_, stream())};
@@ -735,7 +737,7 @@ Let's look at a simple script and its results:
 
     print(f"c shape: {c.shape}")
     print(f"c dtype: {c.dtype}")
-    print(f"c is correct: {mx.all(c == 6.0).item()}")
+    print(f"c correct: {mx.all(c == 6.0).item()}")
 
 Output:
 
@@ -743,7 +745,7 @@ Output:
 
     c shape: [3, 4]
     c dtype: float32
-    c is correct: True
+    c correctness: True
 
 Results
 ^^^^^^^

docs/src/index.rst

@@ -70,7 +70,6 @@ are the CPU and GPU.
    python/fft
    python/linalg
    python/metal
-   python/memory_management
    python/nn
    python/optimizers
    python/distributed

docs/src/python/array.rst

@@ -38,7 +38,6 @@ Array
     array.log10
     array.log1p
     array.log2
-    array.logcumsumexp
     array.logsumexp
     array.max
     array.mean

docs/src/python/linalg.rst

@@ -20,6 +20,5 @@ Linear Algebra
     eigh
     lu
     lu_factor
-    pinv
     solve
     solve_triangular

docs/src/python/memory_management.rst

@@ -1,16 +0,0 @@
-Memory Management
-=================
-
-.. currentmodule:: mlx.core
-
-.. autosummary::
-  :toctree: _autosummary
-
-  get_active_memory
-  get_peak_memory
-  reset_peak_memory
-  get_cache_memory
-  set_memory_limit
-  set_cache_limit
-  set_wired_limit
-  clear_cache

docs/src/python/metal.rst

@@ -8,5 +8,13 @@ Metal
 
   is_available
   device_info
+  get_active_memory
+  get_peak_memory
+  reset_peak_memory
+  get_cache_memory
+  set_memory_limit
+  set_cache_limit
+  set_wired_limit
+  clear_cache
   start_capture
   stop_capture

docs/src/python/ops.rst

@@ -36,12 +36,10 @@ Operations
    bitwise_or
    bitwise_xor
    block_masked_mm
-   broadcast_arrays
    broadcast_to
    ceil
    clip
    concatenate
-   contiguous
    conj
    conjugate
    convolve
@@ -103,7 +101,6 @@ Operations
    log10
    log1p
    logaddexp
-   logcumsumexp
    logical_not
    logical_and
    logical_or

docs/src/python/optimizers/common_optimizers.rst

@@ -18,4 +18,3 @@ Common Optimizers
    AdamW
    Adamax
    Lion
-   MultiOptimizer

docs/src/python/transforms.rst

@@ -9,7 +9,6 @@ Transforms
   :toctree: _autosummary
 
    eval
-   async_eval
    compile
    custom_function
    disable_compile

examples/extensions/axpby/axpby.cpp

@@ -72,7 +72,9 @@ void axpby_impl(
     float alpha_,
     float beta_,
     mx::Stream stream) {
-  out.set_data(mx::allocator::malloc(out.nbytes()));
+  // Allocate the output with `malloc_or_wait` which synchronously allocates
+  // memory, potentially waiting if the system is under memory pressure
+  out.set_data(mx::allocator::malloc_or_wait(out.nbytes()));
 
   // Get the CPU command encoder and register input and output arrays
   auto& encoder = mx::cpu::get_command_encoder(stream);
@@ -158,12 +160,12 @@ void Axpby::eval_gpu(
   // Allocate output memory with strides based on specialization
   if (contiguous_kernel) {
     out.set_data(
-        mx::allocator::malloc(x.data_size() * out.itemsize()),
+        mx::allocator::malloc_or_wait(x.data_size() * out.itemsize()),
         x.data_size(),
         x.strides(),
         x.flags());
   } else {
-    out.set_data(mx::allocator::malloc(out.nbytes()));
+    out.set_data(mx::allocator::malloc_or_wait(out.nbytes()));
   }
 
   // Resolve name of kernel (corresponds to axpby.metal)

mlx/allocator.cpp

@@ -4,11 +4,12 @@
 #include <sstream>
 
 #include "mlx/allocator.h"
+#include "mlx/scheduler.h"
 
 namespace mlx::core::allocator {
 
 Buffer malloc(size_t size) {
-  auto buffer = allocator().malloc(size);
+  auto buffer = allocator().malloc(size, /* allow_swap */ true);
   if (size && !buffer.ptr()) {
     std::ostringstream msg;
     msg << "[malloc] Unable to allocate " << size << " bytes.";
@@ -21,4 +22,45 @@ void free(Buffer buffer) {
   allocator().free(buffer);
 }
 
+Buffer CommonAllocator::malloc(size_t size, bool) {
+  void* ptr = std::malloc(size + sizeof(size_t));
+  if (ptr != nullptr) {
+    *static_cast<size_t*>(ptr) = size;
+  }
+  return Buffer{ptr};
+}
+
+void CommonAllocator::free(Buffer buffer) {
+  std::free(buffer.ptr());
+}
+
+size_t CommonAllocator::size(Buffer buffer) const {
+  if (buffer.ptr() == nullptr) {
+    return 0;
+  }
+  return *static_cast<size_t*>(buffer.ptr());
+}
+
+Buffer malloc_or_wait(size_t size) {
+  auto buffer = allocator().malloc(size);
+
+  while (size && !buffer.ptr() && scheduler::n_active_tasks() > 0) {
+    scheduler::wait_for_one();
+    buffer = allocator().malloc(size);
+  }
+
+  // Try swapping if needed
+  if (size && !buffer.ptr()) {
+    buffer = allocator().malloc(size, /* allow_swap = */ true);
+  }
+
+  if (size && !buffer.ptr()) {
+    std::ostringstream msg;
+    msg << "[malloc_or_wait] Unable to allocate " << size << " bytes.";
+    throw std::runtime_error(msg.str());
+  }
+
+  return buffer;
+}
+
 } // namespace mlx::core::allocator

mlx/allocator.h

@@ -32,10 +32,14 @@ Buffer malloc(size_t size);
 
 void free(Buffer buffer);
 
+// Wait for running tasks to finish and free up memory
+// if allocation fails
+Buffer malloc_or_wait(size_t size);
+
 class Allocator {
   /** Abstract base class for a memory allocator. */
  public:
-  virtual Buffer malloc(size_t size) = 0;
+  virtual Buffer malloc(size_t size, bool allow_swap = false) = 0;
   virtual void free(Buffer buffer) = 0;
   virtual size_t size(Buffer buffer) const = 0;
 
@@ -49,4 +53,16 @@ class Allocator {
 
 Allocator& allocator();
 
+class CommonAllocator : public Allocator {
+  /** A general CPU allocator. */
+ public:
+  virtual Buffer malloc(size_t size, bool allow_swap = false) override;
+  virtual void free(Buffer buffer) override;
+  virtual size_t size(Buffer buffer) const override;
+
+ private:
+  CommonAllocator() = default;
+  friend Allocator& allocator();
+};
+
 } // namespace mlx::core::allocator

mlx/backend/common/CMakeLists.txt

@@ -1,7 +1,6 @@
 target_sources(
   mlx
-  PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/broadcasting.cpp
-          ${CMAKE_CURRENT_SOURCE_DIR}/compiled.cpp
+  PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/compiled.cpp
           ${CMAKE_CURRENT_SOURCE_DIR}/common.cpp
           ${CMAKE_CURRENT_SOURCE_DIR}/load.cpp
           ${CMAKE_CURRENT_SOURCE_DIR}/reduce.cpp

mlx/backend/common/binary.h

@@ -44,14 +44,14 @@ inline void set_binary_op_output_data(
   switch (bopt) {
     case BinaryOpType::ScalarScalar:
       out.set_data(
-          allocator::malloc(out.itemsize()), 1, a.strides(), a.flags());
+          allocator::malloc_or_wait(out.itemsize()), 1, a.strides(), a.flags());
       break;
     case BinaryOpType::ScalarVector:
       if (b_donatable) {
         out.copy_shared_buffer(b);
       } else {
         out.set_data(
-            allocator::malloc(b.data_size() * out.itemsize()),
+            allocator::malloc_or_wait(b.data_size() * out.itemsize()),
             b.data_size(),
             b.strides(),
             b.flags());
@@ -62,7 +62,7 @@ inline void set_binary_op_output_data(
         out.copy_shared_buffer(a);
       } else {
         out.set_data(
-            allocator::malloc(a.data_size() * out.itemsize()),
+            allocator::malloc_or_wait(a.data_size() * out.itemsize()),
             a.data_size(),
             a.strides(),
             a.flags());
@@ -75,7 +75,7 @@ inline void set_binary_op_output_data(
         out.copy_shared_buffer(b);
       } else {
         out.set_data(
-            allocator::malloc(a.data_size() * out.itemsize()),
+            allocator::malloc_or_wait(a.data_size() * out.itemsize()),
             a.data_size(),
             a.strides(),
             a.flags());
@@ -88,7 +88,7 @@ inline void set_binary_op_output_data(
           b_donatable && b.flags().row_contiguous && b.size() == out.size()) {
         out.copy_shared_buffer(b);
       } else {
-        out.set_data(allocator::malloc(out.nbytes()));
+        out.set_data(allocator::malloc_or_wait(out.nbytes()));
       }
       break;
   }

mlx/backend/common/broadcasting.cpp

@@ -1,24 +0,0 @@
-// Copyright © 2024 Apple Inc.
-
-#include "mlx/backend/common/utils.h"
-
-namespace mlx::core {
-
-void broadcast(const array& in, array& out) {
-  if (out.size() == 0) {
-    out.set_data(nullptr);
-    return;
-  }
-  Strides strides(out.ndim(), 0);
-  int diff = out.ndim() - in.ndim();
-  for (int i = in.ndim() - 1; i >= 0; --i) {
-    strides[i + diff] = (in.shape()[i] == 1) ? 0 : in.strides()[i];
-  }
-  auto flags = in.flags();
-  if (out.size() > in.size()) {
-    flags.row_contiguous = flags.col_contiguous = false;
-  }
-  out.copy_shared_buffer(in, strides, flags, in.data_size());
-}
-
-} // namespace mlx::core

mlx/backend/common/broadcasting.h

@@ -1,11 +0,0 @@
-// Copyright © 2024 Apple Inc.
-
-#pragma once
-
-#include "mlx/array.h"
-
-namespace mlx::core {
-
-void broadcast(const array& in, array& out);
-
-} // namespace mlx::core

mlx/backend/common/common.cpp

@@ -1,7 +1,6 @@
 // Copyright © 2024 Apple Inc.
 #include <cassert>
 
-#include "mlx/backend/common/broadcasting.h"
 #include "mlx/backend/common/utils.h"
 #include "mlx/primitives.h"
 
@@ -43,6 +42,23 @@ void AsStrided::eval(const std::vector<array>& inputs, array& out) {
   return out.copy_shared_buffer(in, strides_, flags, data_size, offset_);
 }
 
+void broadcast(const array& in, array& out) {
+  if (out.size() == 0) {
+    out.set_data(nullptr);
+    return;
+  }
+  Strides strides(out.ndim(), 0);
+  int diff = out.ndim() - in.ndim();
+  for (int i = in.ndim() - 1; i >= 0; --i) {
+    strides[i + diff] = (in.shape()[i] == 1) ? 0 : in.strides()[i];
+  }
+  auto flags = in.flags();
+  if (out.size() > in.size()) {
+    flags.row_contiguous = flags.col_contiguous = false;
+  }
+  out.copy_shared_buffer(in, strides, flags, in.data_size());
+}
+
 void Broadcast::eval(const std::vector<array>& inputs, array& out) {
   broadcast(inputs[0], out);
 }
@@ -87,7 +103,7 @@ void ExpandDims::eval(const std::vector<array>& inputs, array& out) {
 
 void NumberOfElements::eval(const std::vector<array>& inputs, array& out) {
   assert(inputs.size() == 1);
-  out.set_data(allocator::malloc(out.nbytes()));
+  out.set_data(allocator::malloc_or_wait(out.nbytes()));
 
   double numel = 1;
   for (auto ax : axes_) {

mlx/backend/common/compiled.cpp

@@ -188,7 +188,7 @@ void compiled_allocate_outputs(
     }
     for (; o < outputs.size(); ++o) {
       outputs[o].set_data(
-          allocator::malloc(data_size * outputs[o].itemsize()),
+          allocator::malloc_or_wait(data_size * outputs[o].itemsize()),
           data_size,
           strides,
           flags);
@@ -211,7 +211,7 @@ void compiled_allocate_outputs(
       }
     }
     for (; o < outputs.size(); ++o) {
-      outputs[o].set_data(allocator::malloc(outputs[o].nbytes()));
+      outputs[o].set_data(allocator::malloc_or_wait(outputs[o].nbytes()));
     }
   }
 }

mlx/backend/common/copy.h

@@ -31,14 +31,14 @@ inline bool set_copy_output_data(const array& in, array& out, CopyType ctype) {
       return true;
     } else {
       out.set_data(
-          allocator::malloc(in.data_size() * out.itemsize()),
+          allocator::malloc_or_wait(in.data_size() * out.itemsize()),
           in.data_size(),
           in.strides(),
           in.flags());
       return false;
     }
   } else {
-    out.set_data(allocator::malloc(out.nbytes()));
+    out.set_data(allocator::malloc_or_wait(out.nbytes()));
     return false;
   }
 }

mlx/backend/common/load.cpp

@@ -28,7 +28,7 @@ void swap_endianness(uint8_t* data_bytes, size_t N) {
 namespace mlx::core {
 
 void Load::eval_cpu(const std::vector<array>& inputs, array& out) {
-  out.set_data(allocator::malloc(out.nbytes()));
+  out.set_data(allocator::malloc_or_wait(out.nbytes()));
   auto read_task = [out_ptr = out.data<char>(),
                     size = out.size(),
                     itemsize = out.itemsize(),

mlx/backend/common/ternary.h

@@ -48,12 +48,12 @@ inline void set_ternary_op_output_data(
   switch (topt) {
     case TernaryOpType::ScalarScalarScalar:
       out.set_data(
-          allocator::malloc(out.itemsize()), 1, b.strides(), b.flags());
+          allocator::malloc_or_wait(out.itemsize()), 1, b.strides(), b.flags());
       break;
     case TernaryOpType::VectorVectorVector:
       if (!(maybe_donate(a) || maybe_donate(b) || maybe_donate(c))) {
         out.set_data(
-            allocator::malloc(out.itemsize() * b.data_size()),
+            allocator::malloc_or_wait(out.itemsize() * b.data_size()),
             b.data_size(),
             b.strides(),
             b.flags());
@@ -64,7 +64,7 @@ inline void set_ternary_op_output_data(
       if (!((a.flags().row_contiguous && maybe_donate(a)) ||
             (b.flags().row_contiguous && maybe_donate(b)) ||
             (c.flags().row_contiguous && maybe_donate(c)))) {
-        out.set_data(allocator::malloc(out.nbytes()));
+        out.set_data(allocator::malloc_or_wait(out.nbytes()));
       }
       break;
   }

mlx/backend/cpu/CMakeLists.txt

@@ -58,7 +58,6 @@ target_sources(
           ${CMAKE_CURRENT_SOURCE_DIR}/scan.cpp
           ${CMAKE_CURRENT_SOURCE_DIR}/select.cpp
           ${CMAKE_CURRENT_SOURCE_DIR}/softmax.cpp
-          ${CMAKE_CURRENT_SOURCE_DIR}/logsumexp.cpp
           ${CMAKE_CURRENT_SOURCE_DIR}/sort.cpp
           ${CMAKE_CURRENT_SOURCE_DIR}/threefry.cpp
           ${CMAKE_CURRENT_SOURCE_DIR}/indexing.cpp
@@ -74,8 +73,8 @@ target_sources(
 if(MLX_BUILD_ACCELERATE)
   target_sources(mlx PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/gemms/bnns.cpp)
 else()
-  target_sources(mlx PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/gemms/simd_fp16.cpp
-                             ${CMAKE_CURRENT_SOURCE_DIR}/gemms/simd_bf16.cpp)
+  target_sources(mlx PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/gemms/no_fp16.cpp
+                             ${CMAKE_CURRENT_SOURCE_DIR}/gemms/no_bf16.cpp)
 endif()
 
 if(IOS)

mlx/backend/cpu/arg_reduce.cpp

@@ -68,7 +68,7 @@ void arg_reduce_dispatch(
 void ArgReduce::eval_cpu(const std::vector<array>& inputs, array& out) {
   assert(inputs.size() == 1);
   auto& in = inputs[0];
-  out.set_data(allocator::malloc(out.nbytes()));
+  out.set_data(allocator::malloc_or_wait(out.nbytes()));
   auto& encoder = cpu::get_command_encoder(stream());
   encoder.set_input_array(in);
   encoder.set_output_array(out);

mlx/backend/cpu/conv.cpp

@@ -921,7 +921,7 @@ void explicit_gemm_conv_1D_cpu(
 
   if (out.dtype() != float32) {
     gemm_out = array(out.shape(), float32, nullptr, {});
-    gemm_out.set_data(allocator::malloc(gemm_out.nbytes()));
+    gemm_out.set_data(allocator::malloc_or_wait(gemm_out.nbytes()));
     temps.push_back(gemm_out);
   }
 
@@ -1048,7 +1048,7 @@ void explicit_gemm_conv_2D_cpu(
 
   if (out.dtype() != float32) {
     gemm_out = array(out.shape(), float32, nullptr, {});
-    gemm_out.set_data(allocator::malloc(gemm_out.nbytes()));
+    gemm_out.set_data(allocator::malloc_or_wait(gemm_out.nbytes()));
     temps.push_back(gemm_out);
   }
 
@@ -1214,7 +1214,7 @@ void explicit_gemm_conv_ND_cpu(
 
   if (out.dtype() != float32) {
     gemm_out = array(out.shape(), float32, nullptr, {});
-    gemm_out.set_data(allocator::malloc(gemm_out.nbytes()));
+    gemm_out.set_data(allocator::malloc_or_wait(gemm_out.nbytes()));
     temps.push_back(gemm_out);
   }
 
@@ -1327,7 +1327,7 @@ void conv_3D_cpu(
 } // namespace
 
 void Convolution::eval_cpu(const std::vector<array>& inputs, array& out) {
-  out.set_data(allocator::malloc(out.nbytes()));
+  out.set_data(allocator::malloc_or_wait(out.nbytes()));
 
   auto& in = inputs[0];
   auto& wt = inputs[1];

mlx/backend/cpu/distributed.cpp

@@ -30,7 +30,7 @@ void AllReduce::eval_cpu(
       if (in.is_donatable()) {
         out.copy_shared_buffer(in);
       } else {
-        out.set_data(allocator::malloc(out.nbytes()));
+        out.set_data(allocator::malloc_or_wait(out.nbytes()));
       }
       return in;
     } else {
@@ -46,15 +46,8 @@ void AllReduce::eval_cpu(
     case Sum:
       distributed::detail::all_sum(group(), in, outputs[0], stream());
       break;
-    case Max:
-      distributed::detail::all_max(group(), in, outputs[0], stream());
-      break;
-    case Min:
-      distributed::detail::all_min(group(), in, outputs[0], stream());
-      break;
     default:
-      throw std::runtime_error(
-          "Only all reduce sum, min and max are supported for now");
+      throw std::runtime_error("Only all reduce sum is supported for now");
   }
 }
 
@@ -65,7 +58,7 @@ void AllGather::eval_cpu(
   assert(outputs.size() == 1);
 
   auto [in, copied] = ensure_row_contiguous(inputs[0], stream());
-  outputs[0].set_data(allocator::malloc(outputs[0].nbytes()));
+  outputs[0].set_data(allocator::malloc_or_wait(outputs[0].nbytes()));
   distributed::detail::all_gather(group(), in, outputs[0], stream());
   if (copied) {
     auto& enc = cpu::get_command_encoder(stream());
@@ -94,7 +87,7 @@ void Recv::eval_cpu(
   assert(inputs.size() == 0);
   assert(outputs.size() == 1);
 
-  outputs[0].set_data(allocator::malloc(outputs[0].nbytes()));
+  outputs[0].set_data(allocator::malloc_or_wait(outputs[0].nbytes()));
   distributed::detail::recv(group(), outputs[0], src_, stream());
 }
 

mlx/backend/cpu/eigh.cpp

@@ -55,8 +55,9 @@ void eigh_impl(
       liwork = iwork;
     }
 
-    auto work_buf = array::Data{allocator::malloc(sizeof(T) * lwork)};
-    auto iwork_buf = array::Data{allocator::malloc(sizeof(int) * liwork)};
+    auto work_buf = array::Data{allocator::malloc_or_wait(sizeof(T) * lwork)};
+    auto iwork_buf =
+        array::Data{allocator::malloc_or_wait(sizeof(int) * liwork)};
     for (size_t i = 0; i < size / (N * N); ++i) {
       syevd<T>(
           &jobz,
@@ -97,7 +98,7 @@ void Eigh::eval_cpu(
       ? outputs[1]
       : array(a.shape(), a.dtype(), nullptr, {});
 
-  values.set_data(allocator::malloc(values.nbytes()));
+  values.set_data(allocator::malloc_or_wait(values.nbytes()));
 
   copy(
       a,

mlx/backend/cpu/encoder.h

@@ -9,9 +9,6 @@
 
 namespace mlx::core::cpu {
 
-// Number of dispatches per scheduler task
-constexpr int DISPATCHES_PER_TASK = 10;
-
 struct CommandEncoder {
   CommandEncoder(Stream stream) : stream_(stream) {}
 
@@ -42,24 +39,13 @@ struct CommandEncoder {
 
   template <class F, class... Args>
   void dispatch(F&& f, Args&&... args) {
-    num_ops_ = (num_ops_ + 1) % DISPATCHES_PER_TASK;
     auto task = std::bind(std::forward<F>(f), std::forward<Args>(args)...);
-    if (num_ops_ == 0) {
-      scheduler::notify_new_task(stream_);
-      auto task_wrap = [s = stream_, task = std::move(task)]() mutable {
-        task();
-        scheduler::notify_task_completion(s);
-      };
-      scheduler::enqueue(stream_, std::move(task_wrap));
-    } else {
-      scheduler::enqueue(stream_, std::move(task));
-    }
+    scheduler::enqueue(stream_, std::move(task));
   }
 
  private:
   Stream stream_;
   std::vector<array> temporaries_;
-  int num_ops_{0};
 };
 
 CommandEncoder& get_command_encoder(Stream stream);

mlx/backend/cpu/eval.cpp

@@ -33,8 +33,12 @@ void eval(array& arr) {
     buffers.erase(it);
   }
   auto& encoder = cpu::get_command_encoder(s);
-  encoder.dispatch([buffers = std::move(buffers),
-                    temps = std::move(encoder.temporaries())]() {});
+  scheduler::notify_new_task(s);
+  encoder.dispatch([s,
+                    buffers = std::move(buffers),
+                    temps = std::move(encoder.temporaries())]() {
+    scheduler::notify_task_completion(s);
+  });
 }
 
 } // namespace mlx::core::cpu

mlx/backend/cpu/fft.cpp

@@ -22,7 +22,7 @@ void FFT::eval_cpu(const std::vector<array>& inputs, array& out) {
     s *= out.itemsize();
   }
 
-  out.set_data(allocator::malloc(out.nbytes()));
+  out.set_data(allocator::malloc_or_wait(out.nbytes()));
 
   std::vector<size_t> shape;
   if (out.dtype() == float32) {

mlx/backend/cpu/gemms/no_bf16.cpp

@@ -0,0 +1,27 @@
+// Copyright © 2025 Apple Inc.
+
+#include "mlx/backend/cpu/gemm.h"
+
+namespace mlx::core {
+
+template <>
+void matmul<bfloat16_t>(
+    const bfloat16_t*,
+    const bfloat16_t*,
+    bfloat16_t*,
+    bool,
+    bool,
+    size_t,
+    size_t,
+    size_t,
+    float,
+    float,
+    size_t,
+    const Shape&,
+    const Strides&,
+    const Shape&,
+    const Strides&) {
+  throw std::runtime_error("[Matmul::eval_cpu] bfloat16 not supported.");
+}
+
+} // namespace mlx::core

mlx/backend/cpu/gemms/no_fp16.cpp

@@ -0,0 +1,27 @@
+// Copyright © 2025 Apple Inc.
+
+#include "mlx/backend/cpu/gemm.h"
+
+namespace mlx::core {
+
+template <>
+void matmul<float16_t>(
+    const float16_t*,
+    const float16_t*,
+    float16_t*,
+    bool,
+    bool,
+    size_t,
+    size_t,
+    size_t,
+    float,
+    float,
+    size_t,
+    const Shape&,
+    const Strides&,
+    const Shape&,
+    const Strides&) {
+  throw std::runtime_error("[Matmul::eval_cpu] float16 not supported.");
+}
+
+} // namespace mlx::core

mlx/backend/cpu/gemms/simd_bf16.cpp

@@ -1,45 +0,0 @@
-// Copyright © 2025 Apple Inc.
-
-#include "mlx/backend/common/utils.h"
-#include "mlx/backend/cpu/gemm.h"
-#include "mlx/backend/cpu/gemms/simd_gemm.h"
-
-namespace mlx::core {
-
-template <>
-void matmul<bfloat16_t>(
-    const bfloat16_t* a,
-    const bfloat16_t* b,
-    bfloat16_t* out,
-    bool a_transposed,
-    bool b_transposed,
-    size_t lda,
-    size_t ldb,
-    size_t ldc,
-    float alpha,
-    float beta,
-    size_t batch_size,
-    const Shape& a_shape,
-    const Strides& a_strides,
-    const Shape& b_shape,
-    const Strides& b_strides) {
-  auto ndim = a_shape.size();
-  size_t M = a_shape[ndim - 2];
-  size_t N = b_shape[ndim - 1];
-  size_t K = a_shape[ndim - 1];
-  for (int i = 0; i < batch_size; ++i) {
-    simd_gemm<bfloat16_t, float>(
-        a + elem_to_loc(M * K * i, a_shape, a_strides),
-        b + elem_to_loc(K * N * i, b_shape, b_strides),
-        out + M * N * i,
-        a_transposed,
-        b_transposed,
-        M,
-        N,
-        K,
-        alpha,
-        beta);
-  }
-}
-
-} // namespace mlx::core

mlx/backend/cpu/gemms/simd_fp16.cpp

@@ -1,45 +0,0 @@
-// Copyright © 2025 Apple Inc.
-
-#include "mlx/backend/common/utils.h"
-#include "mlx/backend/cpu/gemm.h"
-#include "mlx/backend/cpu/gemms/simd_gemm.h"
-
-namespace mlx::core {
-
-template <>
-void matmul<float16_t>(
-    const float16_t* a,
-    const float16_t* b,
-    float16_t* out,
-    bool a_transposed,
-    bool b_transposed,
-    size_t lda,
-    size_t ldb,
-    size_t ldc,
-    float alpha,
-    float beta,
-    size_t batch_size,
-    const Shape& a_shape,
-    const Strides& a_strides,
-    const Shape& b_shape,
-    const Strides& b_strides) {
-  auto ndim = a_shape.size();
-  size_t M = a_shape[ndim - 2];
-  size_t N = b_shape[ndim - 1];
-  size_t K = a_shape[ndim - 1];
-  for (int i = 0; i < batch_size; ++i) {
-    simd_gemm<float16_t, float>(
-        a + elem_to_loc(M * K * i, a_shape, a_strides),
-        b + elem_to_loc(K * N * i, b_shape, b_strides),
-        out + M * N * i,
-        a_transposed,
-        b_transposed,
-        M,
-        N,
-        K,
-        alpha,
-        beta);
-  }
-}
-
-} // namespace mlx::core

mlx/backend/cpu/gemms/simd_gemm.h

@@ -1,139 +0,0 @@
-// Copyright © 2025 Apple Inc.
-#pragma once
-
-#include "mlx/backend/cpu/simd/simd.h"
-
-namespace mlx::core {
-
-inline int ceildiv(int a, int b) {
-  return (a + b - 1) / b;
-}
-
-template <int block_size, typename T, typename AccT>
-void load_block(
-    const T* in,
-    AccT* out,
-    int M,
-    int N,
-    int i,
-    int j,
-    bool transpose) {
-  if (transpose) {
-    for (int ii = 0; ii < block_size && i * block_size + ii < M; ++ii) {
-      for (int jj = 0; jj < block_size && j * block_size + jj < N; ++jj) {
-        out[jj * block_size + ii] =
-            in[(i * block_size + ii) * N + j * block_size + jj];
-      }
-    }
-  } else {
-    for (int ii = 0; ii < block_size && i * block_size + ii < M; ++ii) {
-      for (int jj = 0; jj < block_size && j * block_size + jj < N; ++jj) {
-        out[ii * block_size + jj] =
-            in[(i * block_size + ii) * N + j * block_size + jj];
-      }
-    }
-  }
-}
-
-template <typename T, typename AccT>
-void simd_gemm(
-    const T* a,
-    const T* b,
-    T* c,
-    bool a_trans,
-    bool b_trans,
-    int M,
-    int N,
-    int K,
-    float alpha,
-    float beta) {
-  constexpr int block_size = 16;
-  constexpr int simd_size = simd::max_size<AccT>;
-  static_assert(
-      (block_size % simd_size) == 0,
-      "Block size must be divisible by SIMD size");
-
-  int last_k_block_size = K - block_size * (K / block_size);
-  int last_k_simd_block = (last_k_block_size / simd_size) * simd_size;
-  for (int i = 0; i < ceildiv(M, block_size); i++) {
-    for (int j = 0; j < ceildiv(N, block_size); j++) {
-      AccT c_block[block_size * block_size] = {0.0};
-      AccT a_block[block_size * block_size];
-      AccT b_block[block_size * block_size];
-
-      int k = 0;
-      for (; k < K / block_size; k++) {
-        // Load a and b blocks
-        if (a_trans) {
-          load_block<block_size>(a, a_block, K, M, k, i, true);
-        } else {
-          load_block<block_size>(a, a_block, M, K, i, k, false);
-        }
-        if (b_trans) {
-          load_block<block_size>(b, b_block, N, K, j, k, false);
-        } else {
-          load_block<block_size>(b, b_block, K, N, k, j, true);
-        }
-
-        // Multiply and accumulate
-        for (int ii = 0; ii < block_size && i * block_size + ii < M; ++ii) {
-          for (int jj = 0; jj < block_size && j * block_size + jj < N; ++jj) {
-            for (int kk = 0; kk < block_size; kk += simd_size) {
-              auto av =
-                  simd::load<AccT, simd_size>(a_block + ii * block_size + kk);
-              auto bv =
-                  simd::load<AccT, simd_size>(b_block + jj * block_size + kk);
-              c_block[ii * block_size + jj] += simd::sum(av * bv);
-            }
-          }
-        }
-      }
-      if (last_k_block_size) {
-        // Load a and b blocks
-        if (a_trans) {
-          load_block<block_size>(a, a_block, K, M, k, i, true);
-        } else {
-          load_block<block_size>(a, a_block, M, K, i, k, false);
-        }
-        if (b_trans) {
-          load_block<block_size>(b, b_block, N, K, j, k, false);
-        } else {
-          load_block<block_size>(b, b_block, K, N, k, j, true);
-        }
-
-        // Multiply and accumulate
-        for (int ii = 0; ii < block_size && i * block_size + ii < M; ++ii) {
-          for (int jj = 0; jj < block_size && j * block_size + jj < N; ++jj) {
-            int kk = 0;
-            for (; kk < last_k_simd_block; kk += simd_size) {
-              auto av =
-                  simd::load<AccT, simd_size>(a_block + ii * block_size + kk);
-              auto bv =
-                  simd::load<AccT, simd_size>(b_block + jj * block_size + kk);
-              c_block[ii * block_size + jj] += simd::sum(av * bv);
-            }
-            for (; kk < last_k_block_size; ++kk) {
-              c_block[ii * block_size + jj] +=
-                  a_block[ii * block_size + kk] * b_block[jj * block_size + kk];
-            }
-          }
-        }
-      }
-
-      // Store
-      for (int ii = 0; ii < block_size && i * block_size + ii < M; ++ii) {
-        for (int jj = 0; jj < block_size && j * block_size + jj < N; ++jj) {
-          auto c_idx = (i * block_size + ii) * N + j * block_size + jj;
-          if (beta != 0) {
-            c[c_idx] = static_cast<T>(
-                alpha * c_block[ii * block_size + jj] + beta * c[c_idx]);
-          } else {
-            c[c_idx] = static_cast<T>(alpha * c_block[ii * block_size + jj]);
-          }
-        }
-      }
-    }
-  }
-}
-
-} // namespace mlx::core

mlx/backend/cpu/indexing.cpp

@@ -197,7 +197,7 @@ void dispatch_gather(
 }
 
 void Gather::eval_cpu(const std::vector<array>& inputs, array& out) {
-  out.set_data(allocator::malloc(out.nbytes()));
+  out.set_data(allocator::malloc_or_wait(out.nbytes()));
 
   auto& src = inputs[0];
   std::vector<array> inds;
@@ -354,7 +354,7 @@ void dispatch_gather_axis(
 }
 
 void GatherAxis::eval_cpu(const std::vector<array>& inputs, array& out) {
-  out.set_data(allocator::malloc(out.nbytes()));
+  out.set_data(allocator::malloc_or_wait(out.nbytes()));
 
   auto& src = inputs[0];
   auto& inds = inputs[1];

mlx/backend/cpu/inverse.cpp

@@ -11,7 +11,7 @@ namespace mlx::core {
 template <typename T>
 void general_inv(T* inv, int N) {
   int info;
-  auto ipiv = array::Data{allocator::malloc(sizeof(int) * N)};
+  auto ipiv = array::Data{allocator::malloc_or_wait(sizeof(int) * N)};
   // Compute LU factorization.
   getrf<T>(
       /* m = */ &N,
@@ -49,7 +49,7 @@ void general_inv(T* inv, int N) {
   }
 
   const int lwork = workspace_size;
-  auto scratch = array::Data{allocator::malloc(sizeof(T) * lwork)};
+  auto scratch = array::Data{allocator::malloc_or_wait(sizeof(T) * lwork)};
 
   // Compute inverse.
   getri<T>(

mlx/backend/cpu/logsumexp.cpp

@@ -1,140 +0,0 @@
-// Copyright © 2023-2024 Apple Inc.
-
-#include <cassert>
-#include <cmath>
-
-#include "mlx/backend/cpu/copy.h"
-#include "mlx/backend/cpu/encoder.h"
-#include "mlx/backend/cpu/simd/simd.h"
-#include "mlx/primitives.h"
-#include "mlx/types/limits.h"
-
-namespace mlx::core {
-
-namespace {
-
-using namespace mlx::core::simd;
-
-template <typename T, typename AccT>
-void logsumexp(const array& in, array& out, Stream stream) {
-  auto& encoder = cpu::get_command_encoder(stream);
-  encoder.set_input_array(in);
-  encoder.set_output_array(out);
-
-  const T* in_ptr = in.data<T>();
-  T* out_ptr = out.data<T>();
-
-  int M = in.shape().back();
-  int L = in.data_size() / M;
-
-  encoder.dispatch([in_ptr, out_ptr, M, L]() mutable {
-    constexpr int N = std::min(max_size<AccT>, max_size<T>);
-
-    const T* current_in_ptr;
-
-    for (int i = 0; i < L; i++, in_ptr += M, out_ptr += 1) {
-      // Find the maximum
-      current_in_ptr = in_ptr;
-      Simd<AccT, N> vmaximum(-numeric_limits<AccT>::infinity());
-      size_t s = M;
-      while (s >= N) {
-        Simd<AccT, N> vals = load<T, N>(current_in_ptr);
-        vmaximum = maximum(vals, vmaximum);
-        current_in_ptr += N;
-        s -= N;
-      }
-
-      AccT maximum = max(vmaximum);
-      while (s-- > 0) {
-        maximum = std::max(maximum, static_cast<AccT>(*current_in_ptr));
-        current_in_ptr++;
-      }
-
-      // Compute the normalizer and the exponentials
-      Simd<AccT, N> vnormalizer(0.0);
-      current_in_ptr = in_ptr;
-      s = M;
-      while (s >= N) {
-        Simd<AccT, N> vexp = load<T, N>(current_in_ptr);
-        vexp = exp(vexp - maximum);
-        vnormalizer = vnormalizer + vexp;
-        current_in_ptr += N;
-        s -= N;
-      }
-      AccT normalizer = sum(vnormalizer);
-      while (s-- > 0) {
-        AccT _exp = std::exp(*current_in_ptr - maximum);
-        normalizer += _exp;
-        current_in_ptr++;
-      }
-      // Normalize
-      *out_ptr = std::isinf(maximum)
-          ? static_cast<T>(maximum)
-          : static_cast<T>(std::log(normalizer) + maximum);
-    }
-  });
-}
-
-} // namespace
-
-void LogSumExp::eval_cpu(const std::vector<array>& inputs, array& out) {
-  assert(inputs.size() == 1);
-
-  // Make sure that the last dimension is contiguous
-  auto s = stream();
-  auto& encoder = cpu::get_command_encoder(s);
-  auto ensure_contiguous = [&s, &encoder](const array& x) {
-    if (x.flags().contiguous && x.strides()[x.ndim() - 1] == 1) {
-      return x;
-    } else {
-      auto x_copy = array(x.shape(), x.dtype(), nullptr, {});
-      copy(x, x_copy, CopyType::General, s);
-      encoder.add_temporary(x_copy);
-      return x_copy;
-    }
-  };
-
-  auto in = ensure_contiguous(inputs[0]);
-  if (in.flags().row_contiguous) {
-    out.set_data(allocator::malloc(out.nbytes()));
-  } else {
-    auto n = in.shape(-1);
-    auto flags = in.flags();
-    auto strides = in.strides();
-    for (auto& s : strides) {
-      s /= n;
-    }
-    bool col_contig = strides[0] == 1;
-    for (int i = 1; col_contig && i < strides.size(); ++i) {
-      col_contig &=
-          (out.shape(i) == 1 || strides[i - 1] == out.shape(i) * strides[i]);
-    }
-    flags.col_contiguous = col_contig;
-    out.set_data(
-        allocator::malloc(in.nbytes() / n),
-        in.data_size() / n,
-        std::move(strides),
-        flags);
-  }
-
-  switch (in.dtype()) {
-    case float32:
-      logsumexp<float, float>(in, out, stream());
-      break;
-    case float16:
-      logsumexp<float16_t, float>(in, out, stream());
-      break;
-    case bfloat16:
-      logsumexp<bfloat16_t, float>(in, out, stream());
-      break;
-    case float64:
-      logsumexp<double, double>(in, out, stream());
-      break;
-    default:
-      throw std::runtime_error(
-          "[logsumexp] only supports floating point types");
-      break;
-  }
-}
-
-} // namespace mlx::core

mlx/backend/cpu/luf.cpp

@@ -30,7 +30,8 @@ void luf_impl(
   auto strides = lu.strides();
   strides[ndim - 1] = M;
   strides[ndim - 2] = 1;
-  lu.set_data(allocator::malloc(lu.nbytes()), lu.nbytes(), strides, flags);
+  lu.set_data(
+      allocator::malloc_or_wait(lu.nbytes()), lu.nbytes(), strides, flags);
   copy_inplace(
       a,
       lu,
@@ -43,8 +44,8 @@ void luf_impl(
       stream);
 
   auto a_ptr = lu.data<T>();
-  pivots.set_data(allocator::malloc(pivots.nbytes()));
-  row_indices.set_data(allocator::malloc(row_indices.nbytes()));
+  pivots.set_data(allocator::malloc_or_wait(pivots.nbytes()));
+  row_indices.set_data(allocator::malloc_or_wait(row_indices.nbytes()));
   auto pivots_ptr = pivots.data<uint32_t>();
   auto row_indices_ptr = row_indices.data<uint32_t>();
   size_t num_matrices = a.size() / (M * N);

mlx/backend/cpu/masked_mm.cpp

@@ -59,7 +59,7 @@ void BlockMaskedMM::eval_cpu(const std::vector<array>& inputs, array& out) {
     throw std::runtime_error(
         "[BlockMaskedMM::eval] Currently only supports float32.");
   }
-  out.set_data(allocator::malloc(out.nbytes()));
+  out.set_data(allocator::malloc_or_wait(out.nbytes()));
 
   auto& a_pre = inputs[0];
   auto& b_pre = inputs[1];
@@ -318,7 +318,7 @@ void GatherMM::eval_cpu(const std::vector<array>& inputs, array& out) {
     throw std::runtime_error(
         "[GatherMM::eval] Currently only supports float32.");
   }
-  out.set_data(allocator::malloc(out.nbytes()));
+  out.set_data(allocator::malloc_or_wait(out.nbytes()));
 
   auto& a_pre = inputs[0];
   auto& b_pre = inputs[1];

mlx/backend/cpu/matmul.cpp

@@ -115,7 +115,7 @@ void matmul_general(
 }
 
 void Matmul::eval_cpu(const std::vector<array>& inputs, array& out) {
-  out.set_data(allocator::malloc(out.nbytes()));
+  out.set_data(allocator::malloc_or_wait(out.nbytes()));
   if (inputs[0].shape(-1) == 0) {
     auto& encoder = cpu::get_command_encoder(stream());
     encoder.set_output_array(out);

mlx/backend/cpu/primitives.cpp

@@ -21,7 +21,7 @@ namespace mlx::core {
 void reshape(const array& in, array& out) {
   auto [copy_necessary, out_strides] = prepare_reshape(in, out);
   if (copy_necessary) {
-    out.set_data(allocator::malloc(out.nbytes()));
+    out.set_data(allocator::malloc_or_wait(out.nbytes()));
     copy_inplace(in, out, CopyType::General, out.primitive().stream());
   } else {
     shared_buffer_reshape(in, out_strides, out);
@@ -39,7 +39,7 @@ static std::pair<array, bool> compute_dynamic_offset(
   if (donate) {
     offset.copy_shared_buffer(indices);
   } else {
-    offset.set_data(allocator::malloc(offset.itemsize()));
+    offset.set_data(allocator::malloc_or_wait(offset.itemsize()));
   }
 
   auto& encoder = cpu::get_command_encoder(stream);
@@ -124,7 +124,7 @@ void Transpose::eval_cpu(const std::vector<array>& inputs, array& out) {
 
 void Arange::eval_cpu(const std::vector<array>& inputs, array& out) {
   assert(inputs.size() == 0);
-  out.set_data(allocator::malloc(out.nbytes()));
+  out.set_data(allocator::malloc_or_wait(out.nbytes()));
   switch (out.dtype()) {
     case bool_:
       throw std::runtime_error("Bool type unsupported for arange.");
@@ -186,7 +186,7 @@ void Concatenate::eval_cpu(const std::vector<array>& inputs, array& out) {
   }
   std::partial_sum(sizes.cbegin(), sizes.cend(), sizes.begin());
 
-  out.set_data(allocator::malloc(out.nbytes()));
+  out.set_data(allocator::malloc_or_wait(out.nbytes()));
 
   auto strides = out.strides();
   auto flags = out.flags();
@@ -205,10 +205,8 @@ void Concatenate::eval_cpu(const std::vector<array>& inputs, array& out) {
 void Contiguous::eval_cpu(const std::vector<array>& inputs, array& out) {
   assert(inputs.size() == 1);
   auto& in = inputs[0];
-  constexpr size_t extra_bytes = 16384;
-  if (in.buffer_size() <= out.nbytes() + extra_bytes &&
-      (in.flags().row_contiguous ||
-       (allow_col_major_ && in.flags().col_contiguous))) {
+  if (in.flags().row_contiguous ||
+      (allow_col_major_ && in.flags().col_contiguous)) {
     out.copy_shared_buffer(in);
   } else {
     copy(in, out, CopyType::General, stream());
@@ -278,7 +276,7 @@ void RandomBits::eval_cpu(const std::vector<array>& inputs, array& out) {
 
   size_t elems_per_key = out.size() / num_keys;
   size_t bytes_per_key = out.itemsize() * elems_per_key;
-  out.set_data(allocator::malloc(out.nbytes()));
+  out.set_data(allocator::malloc_or_wait(out.nbytes()));
 
   auto kptr = inputs[0].data<uint32_t>();
   auto cptr = out.data<char>();
@@ -337,7 +335,7 @@ void DynamicSlice::eval_cpu(const std::vector<array>& inputs, array& out) {
     return;
   }
   auto& in = inputs[0];
-  out.set_data(allocator::malloc(out.nbytes()));
+  out.set_data(allocator::malloc_or_wait(out.nbytes()));
   auto [in_offset, donated] =
       compute_dynamic_offset(inputs[1], in.strides(), axes_, stream());
   copy_inplace(
@@ -452,7 +450,7 @@ void View::eval_cpu(const std::vector<array>& inputs, array& out) {
   } else {
     auto tmp = array(
         in.shape(), in.dtype() == bool_ ? uint8 : in.dtype(), nullptr, {});
-    tmp.set_data(allocator::malloc(tmp.nbytes()));
+    tmp.set_data(allocator::malloc_or_wait(tmp.nbytes()));
     if (in.dtype() == bool_) {
       auto in_tmp = array(in.shape(), uint8, nullptr, {});
       in_tmp.copy_shared_buffer(in);

mlx/backend/cpu/qrf.cpp

@@ -25,11 +25,12 @@ void qrf_impl(const array& a, array& q, array& r, Stream stream) {
   auto strides = in.strides();
   strides[in.ndim() - 2] = 1;
   strides[in.ndim() - 1] = M;
-  in.set_data(allocator::malloc(in.nbytes()), in.nbytes(), strides, flags);
+  in.set_data(
+      allocator::malloc_or_wait(in.nbytes()), in.nbytes(), strides, flags);
   copy_inplace(a, in, CopyType::GeneralGeneral, stream);
   auto& encoder = cpu::get_command_encoder(stream);
-  q.set_data(allocator::malloc(q.nbytes()));
-  r.set_data(allocator::malloc(r.nbytes()));
+  q.set_data(allocator::malloc_or_wait(q.nbytes()));
+  r.set_data(allocator::malloc_or_wait(r.nbytes()));
 
   auto in_ptr = in.data<T>();
   auto r_ptr = r.data<T>();
@@ -40,7 +41,8 @@ void qrf_impl(const array& a, array& q, array& r, Stream stream) {
   encoder.set_output_array(r);
   encoder.dispatch([in_ptr, q_ptr, r_ptr, M, N, lda, num_matrices]() {
     int num_reflectors = std::min(M, N);
-    auto tau = allocator::malloc(sizeof(T) * num_matrices * num_reflectors);
+    auto tau =
+        allocator::malloc_or_wait(sizeof(T) * num_matrices * num_reflectors);
 
     T optimal_work;
     int lwork = -1;
@@ -51,7 +53,7 @@ void qrf_impl(const array& a, array& q, array& r, Stream stream) {
 
     // Update workspace size
     lwork = optimal_work;
-    auto work = allocator::malloc(sizeof(T) * lwork);
+    auto work = allocator::malloc_or_wait(sizeof(T) * lwork);
 
     // Loop over matrices
     for (int i = 0; i < num_matrices; ++i) {
@@ -94,7 +96,7 @@ void qrf_impl(const array& a, array& q, array& r, Stream stream) {
         &lwork,
         &info);
     lwork = optimal_work;
-    work = allocator::malloc(sizeof(T) * lwork);
+    work = allocator::malloc_or_wait(sizeof(T) * lwork);
 
     // Loop over matrices
     for (int i = 0; i < num_matrices; ++i) {

mlx/backend/cpu/quantized.cpp

@@ -515,7 +515,7 @@ void QuantizedMatmul::eval_cpu(const std::vector<array>& inputs, array& out) {
   auto scales = ensure_row_contiguous(scales_pre);
   auto biases = ensure_row_contiguous(biases_pre);
 
-  out.set_data(allocator::malloc(out.nbytes()));
+  out.set_data(allocator::malloc_or_wait(out.nbytes()));
 
   auto& encoder = cpu::get_command_encoder(stream());
   encoder.add_temporaries(std::move(temps));
@@ -565,7 +565,7 @@ void GatherQMM::eval_cpu(const std::vector<array>& inputs, array& out) {
   auto scales = ensure_row_contiguous_last_dims(scales_pre);
   auto biases = ensure_row_contiguous_last_dims(biases_pre);
 
-  out.set_data(allocator::malloc(out.nbytes()));
+  out.set_data(allocator::malloc_or_wait(out.nbytes()));
 
   auto& encoder = cpu::get_command_encoder(stream());
   encoder.add_temporaries(std::move(temps));
@@ -691,12 +691,12 @@ void fast::AffineQuantize::eval_cpu(
 
   auto [w, copied] = ensure_row_contiguous(inputs[0]);
   auto& out = outputs[0];
-  out.set_data(allocator::malloc(out.nbytes()));
+  out.set_data(allocator::malloc_or_wait(out.nbytes()));
 
   auto& scales = outputs[1];
   auto& biases = outputs[2];
-  scales.set_data(allocator::malloc(scales.nbytes()));
-  biases.set_data(allocator::malloc(biases.nbytes()));
+  scales.set_data(allocator::malloc_or_wait(scales.nbytes()));
+  biases.set_data(allocator::malloc_or_wait(biases.nbytes()));
   auto& encoder = cpu::get_command_encoder(stream());
   if (copied) {
     encoder.add_temporary(w);

mlx/backend/cpu/reduce.cpp

@@ -433,7 +433,7 @@ void reduce_dispatch_min_max(
 void Reduce::eval_cpu(const std::vector<array>& inputs, array& out) {
   assert(inputs.size() == 1);
   auto& in = inputs[0];
-  out.set_data(allocator::malloc(out.nbytes()));
+  out.set_data(allocator::malloc_or_wait(out.nbytes()));
   auto& encoder = cpu::get_command_encoder(stream());
   encoder.set_input_array(in);
   encoder.set_output_array(out);

mlx/backend/cpu/scan.cpp

@@ -3,7 +3,6 @@
 #include <cassert>
 
 #include "mlx/backend/common/utils.h"
-#include "mlx/backend/cpu/binary_ops.h"
 #include "mlx/backend/cpu/copy.h"
 #include "mlx/backend/cpu/encoder.h"
 #include "mlx/backend/cpu/simd/simd.h"
@@ -227,16 +226,6 @@ void scan_dispatch(
       scan_op<T, U>(in, out, axis, reverse, inclusive, op, init);
       break;
     }
-    case Scan::LogAddExp: {
-      auto op = [](U a, T b) {
-        return detail::LogAddExp{}(a, static_cast<U>(b));
-      };
-      auto init = (issubdtype(in.dtype(), floating))
-          ? static_cast<U>(-std::numeric_limits<float>::infinity())
-          : std::numeric_limits<U>::min();
-      scan_op<T, U>(in, out, axis, reverse, inclusive, op, init);
-      break;
-    }
   }
 }
 
@@ -255,7 +244,7 @@ void Scan::eval_cpu(const std::vector<array>& inputs, array& out) {
     in = arr_copy;
     encoder.add_temporary(arr_copy);
   }
-  out.set_data(allocator::malloc(out.nbytes()));
+  out.set_data(allocator::malloc_or_wait(out.nbytes()));
 
   encoder.set_input_array(in);
   encoder.set_output_array(out);

mlx/backend/cpu/simd/accelerate_fp16_simd.h

@@ -17,7 +17,7 @@ struct ScalarT<float16_t, N> {
 #endif
 
 template <>
-inline constexpr int max_size<float16_t> = N;
+static constexpr int max_size<float16_t> = N;
 
 #define SIMD_FP16_DEFAULT_UNARY(op)                    \
   template <>                                          \

mlx/backend/cpu/simd/accelerate_simd.h

@@ -83,25 +83,25 @@ struct Simd {
 // Values chosen based on benchmarks on M3 Max
 // TODO: consider choosing these more optimally
 template <>
-inline constexpr int max_size<int8_t> = 16;
+static constexpr int max_size<int8_t> = 16;
 template <>
-inline constexpr int max_size<int16_t> = 16;
+static constexpr int max_size<int16_t> = 16;
 template <>
-inline constexpr int max_size<int> = 8;
+static constexpr int max_size<int> = 8;
 template <>
-inline constexpr int max_size<int64_t> = 4;
+static constexpr int max_size<int64_t> = 4;
 template <>
-inline constexpr int max_size<uint8_t> = 16;
+static constexpr int max_size<uint8_t> = 16;
 template <>
-inline constexpr int max_size<uint16_t> = 16;
+static constexpr int max_size<uint16_t> = 16;
 template <>
-inline constexpr int max_size<uint32_t> = 8;
+static constexpr int max_size<uint32_t> = 8;
 template <>
-inline constexpr int max_size<uint64_t> = 4;
+static constexpr int max_size<uint64_t> = 4;
 template <>
-inline constexpr int max_size<float> = 8;
+static constexpr int max_size<float> = 8;
 template <>
-inline constexpr int max_size<double> = 4;
+static constexpr int max_size<double> = 4;
 
 #define SIMD_DEFAULT_UNARY(name, op) \
   template <typename T, int N>       \

mlx/backend/cpu/simd/base_simd.h

@@ -87,6 +87,7 @@ DEFAULT_UNARY(cosh, std::cosh)
 DEFAULT_UNARY(expm1, std::expm1)
 DEFAULT_UNARY(floor, std::floor)
 DEFAULT_UNARY(log, std::log)
+DEFAULT_UNARY(log2, std::log2)
 DEFAULT_UNARY(log10, std::log10)
 DEFAULT_UNARY(log1p, std::log1p)
 DEFAULT_UNARY(sinh, std::sinh)
@@ -94,17 +95,6 @@ DEFAULT_UNARY(sqrt, std::sqrt)
 DEFAULT_UNARY(tan, std::tan)
 DEFAULT_UNARY(tanh, std::tanh)
 
-template <typename T>
-Simd<T, 1> log2(Simd<T, 1> in) {
-  if constexpr (is_complex<T>) {
-    auto out = std::log(in.value);
-    auto scale = decltype(out.real())(M_LN2);
-    return Simd<T, 1>{T{out.real() / scale, out.imag() / scale}};
-  } else {
-    return Simd<T, 1>{std::log2(in.value)};
-  }
-}
-
 template <typename T>
 Simd<T, 1> operator~(Simd<T, 1> in) {
   return ~in.value;

mlx/backend/cpu/softmax.cpp

@@ -119,12 +119,17 @@ void Softmax::eval_cpu(const std::vector<array>& inputs, array& out) {
 
   // Make sure that the last dimension is contiguous
   auto set_output = [s = stream(), &out](const array& x) {
-    if (x.flags().contiguous && x.strides()[x.ndim() - 1] == 1) {
+    bool no_copy = x.strides()[x.ndim() - 1] == 1;
+    if (x.ndim() > 1) {
+      auto s = x.strides()[x.ndim() - 2];
+      no_copy &= (s == 0 || s == x.shape().back());
+    }
+    if (no_copy) {
       if (x.is_donatable()) {
         out.copy_shared_buffer(x);
       } else {
         out.set_data(
-            allocator::malloc(x.data_size() * x.itemsize()),
+            allocator::malloc_or_wait(x.data_size() * x.itemsize()),
             x.data_size(),
             x.strides(),
             x.flags());
@@ -141,6 +146,18 @@ void Softmax::eval_cpu(const std::vector<array>& inputs, array& out) {
   auto in = set_output(inputs[0]);
 
   switch (in.dtype()) {
+    case bool_:
+    case uint8:
+    case uint16:
+    case uint32:
+    case uint64:
+    case int8:
+    case int16:
+    case int32:
+    case int64:
+      throw std::runtime_error(
+          "Softmax is defined only for floating point types");
+      break;
     case float32:
       softmax<float, float>(in, out, stream());
       break;
@@ -161,9 +178,9 @@ void Softmax::eval_cpu(const std::vector<array>& inputs, array& out) {
     case float64:
       softmax<double, double>(in, out, stream());
       break;
-    default:
-      throw std::runtime_error(
-          "[softmax] Only defined for floating point types.");
+    case complex64:
+      throw std::invalid_argument(
+          "[Softmax] Not yet implemented for complex64");
       break;
   }
 }

mlx/backend/cpu/sort.cpp

@@ -288,7 +288,7 @@ void ArgSort::eval_cpu(const std::vector<array>& inputs, array& out) {
   auto& in = inputs[0];
 
   // Allocate output
-  out.set_data(allocator::malloc(out.nbytes()));
+  out.set_data(allocator::malloc_or_wait(out.nbytes()));
 
   auto& encoder = cpu::get_command_encoder(stream());
   encoder.set_input_array(in);
@@ -379,7 +379,7 @@ void ArgPartition::eval_cpu(const std::vector<array>& inputs, array& out) {
   auto& in = inputs[0];
 
   // Allocate output
-  out.set_data(allocator::malloc(out.nbytes()));
+  out.set_data(allocator::malloc_or_wait(out.nbytes()));
 
   auto& encoder = cpu::get_command_encoder(stream());
   encoder.set_input_array(in);

mlx/backend/cpu/svd.cpp

@@ -50,9 +50,9 @@ void svd_impl(
     array& s = outputs[1];
     array& vt = outputs[2];
 
-    u.set_data(allocator::malloc(u.nbytes()));
-    s.set_data(allocator::malloc(s.nbytes()));
-    vt.set_data(allocator::malloc(vt.nbytes()));
+    u.set_data(allocator::malloc_or_wait(u.nbytes()));
+    s.set_data(allocator::malloc_or_wait(s.nbytes()));
+    vt.set_data(allocator::malloc_or_wait(vt.nbytes()));
 
     encoder.set_output_array(u);
     encoder.set_output_array(s);
@@ -64,7 +64,7 @@ void svd_impl(
   } else {
     array& s = outputs[0];
 
-    s.set_data(allocator::malloc(s.nbytes()));
+    s.set_data(allocator::malloc_or_wait(s.nbytes()));
 
     encoder.set_output_array(s);
 
@@ -91,7 +91,7 @@ void svd_impl(
 
     // Will contain the indices of eigenvectors that failed to converge (not
     // used here but required by lapack).
-    auto iwork = array::Data{allocator::malloc(sizeof(int) * 12 * K)};
+    auto iwork = array::Data{allocator::malloc_or_wait(sizeof(int) * 12 * K)};
 
     static const int lwork_query = -1;
 
@@ -132,7 +132,7 @@ void svd_impl(
     }
 
     const int lwork = workspace_dimension;
-    auto scratch = array::Data{allocator::malloc(sizeof(T) * lwork)};
+    auto scratch = array::Data{allocator::malloc_or_wait(sizeof(T) * lwork)};
 
     // Loop over matrices.
     for (int i = 0; i < num_matrices; i++) {

mlx/backend/cpu/unary.cpp

@@ -1,8 +1,5 @@
 // Copyright © 2024 Apple Inc.
 
-// Required for using M_LN2 in MSVC.
-#define _USE_MATH_DEFINES
-
 #include <cassert>
 
 #include "mlx/backend/cpu/unary.h"

mlx/backend/cpu/unary.h

@@ -18,13 +18,13 @@ void set_unary_output_data(const array& in, array& out) {
     } else {
       auto size = in.data_size();
       out.set_data(
-          allocator::malloc(size * out.itemsize()),
+          allocator::malloc_or_wait(size * out.itemsize()),
           size,
           in.strides(),
           in.flags());
     }
   } else {
-    out.set_data(allocator::malloc(out.nbytes()));
+    out.set_data(allocator::malloc_or_wait(out.nbytes()));
   }
 }
 

mlx/backend/cpu/unary_ops.h

@@ -86,14 +86,13 @@ struct Sign {
   template <int N, typename T>
   Simd<T, N> operator()(Simd<T, N> x) {
     auto z = Simd<T, N>{0};
-    auto o = Simd<T, N>{1};
-    auto m = Simd<T, N>{-1};
     if constexpr (std::is_unsigned_v<T>) {
-      return simd::select(x == z, z, o);
+      return x != z;
     } else if constexpr (std::is_same_v<T, complex64_t>) {
       return simd::select(x == z, x, Simd<T, N>(x / simd::abs(x)));
     } else {
-      return simd::select(x < z, m, simd::select(x > z, o, z));
+      return simd::select(
+          x < z, Simd<T, N>{-1}, simd::select(x > z, Simd<T, N>{1}, z));
     }
   }
   SINGLE()

mlx/backend/metal/CMakeLists.txt

@@ -47,7 +47,6 @@ if(MLX_METAL_JIT)
   make_jit_source(binary)
   make_jit_source(binary_two)
   make_jit_source(fft kernels/fft/radix.h kernels/fft/readwrite.h)
-  make_jit_source(logsumexp)
   make_jit_source(ternary)
   make_jit_source(softmax)
   make_jit_source(scan)
@@ -61,7 +60,6 @@ if(MLX_METAL_JIT)
     kernels/steel/gemm/transforms.h)
   make_jit_source(steel/gemm/kernels/steel_gemm_fused)
   make_jit_source(steel/gemm/kernels/steel_gemm_masked kernels/steel/defines.h)
-  make_jit_source(steel/gemm/kernels/steel_gemm_gather)
   make_jit_source(steel/gemm/kernels/steel_gemm_splitk)
   make_jit_source(
     steel/conv/conv
@@ -97,7 +95,6 @@ target_sources(
           ${CMAKE_CURRENT_SOURCE_DIR}/fft.cpp
           ${CMAKE_CURRENT_SOURCE_DIR}/hadamard.cpp
           ${CMAKE_CURRENT_SOURCE_DIR}/indexing.cpp
-          ${CMAKE_CURRENT_SOURCE_DIR}/logsumexp.cpp
           ${CMAKE_CURRENT_SOURCE_DIR}/matmul.cpp
           ${CMAKE_CURRENT_SOURCE_DIR}/scaled_dot_product_attention.cpp
           ${CMAKE_CURRENT_SOURCE_DIR}/metal.cpp

mlx/backend/metal/allocator.cpp

@@ -3,7 +3,6 @@
 #include "mlx/backend/metal/metal.h"
 #include "mlx/backend/metal/metal_impl.h"
 #include "mlx/backend/metal/resident.h"
-#include "mlx/memory.h"
 
 #include <mach/vm_page_size.h>
 #include <unistd.h>
@@ -21,9 +20,6 @@ Allocator& allocator() {
 }
 
 void* Buffer::raw_ptr() {
-  if (!ptr_) {
-    return nullptr;
-  }
   return static_cast<MTL::Buffer*>(ptr_)->contents();
 }
 
@@ -33,11 +29,8 @@ namespace metal {
 
 namespace {
 
-BufferCache::BufferCache(ResidencySet& residency_set)
-    : head_(nullptr),
-      tail_(nullptr),
-      pool_size_(0),
-      residency_set_(residency_set) {}
+BufferCache::BufferCache(MTL::Device* device)
+    : device_(device), head_(nullptr), tail_(nullptr), pool_size_(0) {}
 
 BufferCache::~BufferCache() {
   auto pool = metal::new_scoped_memory_pool();
@@ -48,9 +41,6 @@ int BufferCache::clear() {
   int n_release = 0;
   for (auto& [size, holder] : buffer_pool_) {
     if (holder->buf) {
-      if (!holder->buf->heap()) {
-        residency_set_.erase(holder->buf);
-      }
       holder->buf->release();
       n_release++;
     }
@@ -108,9 +98,6 @@ int BufferCache::release_cached_buffers(size_t min_bytes_to_free) {
     while (tail_ && (total_bytes_freed < min_bytes_to_free)) {
       if (tail_->buf) {
         total_bytes_freed += tail_->buf->length();
-        if (!tail_->buf->heap()) {
-          residency_set_.erase(tail_->buf);
-        }
         tail_->buf->release();
         tail_->buf = nullptr;
         n_release++;
@@ -165,7 +152,7 @@ void BufferCache::remove_from_list(BufferCache::BufferHolder* to_remove) {
 MetalAllocator::MetalAllocator()
     : device_(device(mlx::core::Device::gpu).mtl_device()),
       residency_set_(device_),
-      buffer_cache_(residency_set_) {
+      buffer_cache_(device_) {
   auto pool = metal::new_scoped_memory_pool();
   auto memsize = std::get<size_t>(device_info().at("memory_size"));
   auto max_rec_size =
@@ -202,19 +189,16 @@ size_t MetalAllocator::set_cache_limit(size_t limit) {
   return limit;
 };
 
-size_t MetalAllocator::set_memory_limit(size_t limit) {
+size_t MetalAllocator::set_memory_limit(size_t limit, bool relaxed) {
   std::unique_lock lk(mutex_);
   std::swap(limit, block_limit_);
+  relaxed_ = relaxed;
   gc_limit_ = std::min(
       block_limit_,
       static_cast<size_t>(0.95 * device_->recommendedMaxWorkingSetSize()));
   return limit;
 };
 
-size_t MetalAllocator::get_memory_limit() {
-  return block_limit_;
-}
-
 size_t MetalAllocator::set_wired_limit(size_t limit) {
   std::unique_lock lk(mutex_);
   std::swap(limit, wired_limit_);
@@ -222,7 +206,7 @@ size_t MetalAllocator::set_wired_limit(size_t limit) {
   return limit;
 };
 
-Buffer MetalAllocator::malloc(size_t size) {
+Buffer MetalAllocator::malloc(size_t size, bool allow_swap /* = false */) {
   // Metal doesn't like empty buffers
   if (size == 0) {
     return Buffer{nullptr};
@@ -249,6 +233,11 @@ Buffer MetalAllocator::malloc(size_t size) {
   if (!buf) {
     size_t mem_required = get_active_memory() + get_cache_memory() + size;
 
+    // If there is too much memory pressure, fail (likely causes a wait).
+    if (!(allow_swap && relaxed_) && mem_required >= block_limit_) {
+      return Buffer{nullptr};
+    }
+
     auto pool = metal::new_scoped_memory_pool();
 
     // If we have a lot of memory pressure or are over the maximum cache size,
@@ -272,13 +261,9 @@ Buffer MetalAllocator::malloc(size_t size) {
     if (!buf) {
       buf = device_->newBuffer(size, resource_options);
     }
-    if (!buf) {
-      return Buffer{nullptr};
-    }
     lk.lock();
-    num_resources_++;
-    if (!buf->heap()) {
-      residency_set_.insert(buf);
+    if (buf) {
+      num_resources_++;
     }
   }
 
@@ -292,6 +277,10 @@ Buffer MetalAllocator::malloc(size_t size) {
         get_cache_memory() - max_pool_size_);
   }
 
+  if (!buf->heap()) {
+    residency_set_.insert(buf);
+  }
+
   return Buffer{static_cast<void*>(buf)};
 }
 
@@ -307,14 +296,14 @@ void MetalAllocator::free(Buffer buffer) {
     return;
   }
   std::unique_lock lk(mutex_);
+  if (!buf->heap()) {
+    residency_set_.erase(buf);
+  }
   active_memory_ -= buf->length();
   if (get_cache_memory() < max_pool_size_) {
     buffer_cache_.recycle_to_cache(buf);
   } else {
     num_resources_--;
-    if (!buf->heap()) {
-      residency_set_.erase(buf);
-    }
     lk.unlock();
     auto pool = metal::new_scoped_memory_pool();
     buf->release();
@@ -333,40 +322,37 @@ MetalAllocator& allocator() {
   return *allocator_;
 }
 
-} // namespace metal
-
 size_t set_cache_limit(size_t limit) {
-  return metal::allocator().set_cache_limit(limit);
+  return allocator().set_cache_limit(limit);
 }
-size_t set_memory_limit(size_t limit) {
-  return metal::allocator().set_memory_limit(limit);
-}
-size_t get_memory_limit() {
-  return metal::allocator().get_memory_limit();
+size_t set_memory_limit(size_t limit, bool relaxed /* = true */) {
+  return allocator().set_memory_limit(limit, relaxed);
 }
 size_t set_wired_limit(size_t limit) {
-  if (limit > std::get<size_t>(metal::device_info().at(
-                  "max_recommended_working_set_size"))) {
+  if (limit >
+      std::get<size_t>(device_info().at("max_recommended_working_set_size"))) {
     throw std::invalid_argument(
         "[metal::set_wired_limit] Setting a wired limit larger than "
         "the maximum working set size is not allowed.");
   }
-  return metal::allocator().set_wired_limit(limit);
+  return allocator().set_wired_limit(limit);
 }
 size_t get_active_memory() {
-  return metal::allocator().get_active_memory();
+  return allocator().get_active_memory();
 }
 size_t get_peak_memory() {
-  return metal::allocator().get_peak_memory();
+  return allocator().get_peak_memory();
 }
 void reset_peak_memory() {
-  metal::allocator().reset_peak_memory();
+  allocator().reset_peak_memory();
 }
 size_t get_cache_memory() {
-  return metal::allocator().get_cache_memory();
+  return allocator().get_cache_memory();
 }
 void clear_cache() {
-  return metal::allocator().clear_cache();
+  return allocator().clear_cache();
 }
 
+} // namespace metal
+
 } // namespace mlx::core

mlx/backend/metal/allocator.h

@@ -18,7 +18,7 @@ namespace {
 
 class BufferCache {
  public:
-  BufferCache(ResidencySet& residency_set);
+  BufferCache(MTL::Device* device);
   ~BufferCache();
 
   MTL::Buffer* reuse_from_cache(size_t size);
@@ -42,11 +42,13 @@ class BufferCache {
   void add_at_head(BufferHolder* to_add);
   void remove_from_list(BufferHolder* to_remove);
 
+  MTL::Device* device_;
+  MTL::Heap* heap_{nullptr};
+
   std::multimap<size_t, BufferHolder*> buffer_pool_;
   BufferHolder* head_;
   BufferHolder* tail_;
   size_t pool_size_;
-  ResidencySet& residency_set_;
 };
 
 } // namespace
@@ -54,7 +56,7 @@ class BufferCache {
 class MetalAllocator : public allocator::Allocator {
   /** Allocator for Metal GPUs. */
  public:
-  virtual Buffer malloc(size_t size) override;
+  virtual Buffer malloc(size_t size, bool allow_swap = false) override;
   virtual void free(Buffer buffer) override;
   virtual size_t size(Buffer buffer) const override;
   size_t get_active_memory() {
@@ -71,8 +73,7 @@ class MetalAllocator : public allocator::Allocator {
     return buffer_cache_.cache_size();
   };
   size_t set_cache_limit(size_t limit);
-  size_t set_memory_limit(size_t limit);
-  size_t get_memory_limit();
+  size_t set_memory_limit(size_t limit, bool relaxed);
   size_t set_wired_limit(size_t limit);
   void clear_cache();
 

mlx/backend/metal/conv.cpp

@@ -37,7 +37,7 @@ void explicit_gemm_conv_ND_gpu(
   Shape unfolded_shape{implicit_M, implicit_K};
   array in_unfolded(unfolded_shape, in.dtype(), nullptr, {});
 
-  in_unfolded.set_data(allocator::malloc(in_unfolded.nbytes()));
+  in_unfolded.set_data(allocator::malloc_or_wait(in_unfolded.nbytes()));
 
   // Prepare unfolding kernel
   std::ostringstream kname;
@@ -115,7 +115,7 @@ void explicit_gemm_conv_group_ND_gpu(
   // Prepare unfolding array
   Shape unfolded_shape{implicit_M, implicit_K * groups};
   array in_unfolded(unfolded_shape, in.dtype(), nullptr, {});
-  in_unfolded.set_data(allocator::malloc(in_unfolded.nbytes()));
+  in_unfolded.set_data(allocator::malloc_or_wait(in_unfolded.nbytes()));
 
   // Prepare unfolding kernel
   std::ostringstream kname;
@@ -613,7 +613,7 @@ void winograd_conv_2D_gpu(
   // Do filter transform
   Shape filt_wg_shape = {8 * 8, conv_params.C, conv_params.O};
   array filt_wg(std::move(filt_wg_shape), wt.dtype(), nullptr, {});
-  filt_wg.set_data(allocator::malloc(filt_wg.nbytes()));
+  filt_wg.set_data(allocator::malloc_or_wait(filt_wg.nbytes()));
   copies_w.push_back(filt_wg);
   {
     int bc = 32;
@@ -640,7 +640,7 @@ void winograd_conv_2D_gpu(
   // Do input transform
   Shape inp_wg_shape = {8 * 8, N_tiles, conv_params.C};
   array inp_wg(std::move(inp_wg_shape), in.dtype(), nullptr, {});
-  inp_wg.set_data(allocator::malloc(inp_wg.nbytes()));
+  inp_wg.set_data(allocator::malloc_or_wait(inp_wg.nbytes()));
   copies_w.push_back(inp_wg);
   {
     int bc = 32;
@@ -667,7 +667,7 @@ void winograd_conv_2D_gpu(
   // Do batched gemm
   Shape out_wg_shape = {8 * 8, N_tiles, conv_params.O};
   array out_wg(std::move(out_wg_shape), in.dtype(), nullptr, {});
-  out_wg.set_data(allocator::malloc(out_wg.nbytes()));
+  out_wg.set_data(allocator::malloc_or_wait(out_wg.nbytes()));
   copies_w.push_back(out_wg);
   {
     std::vector<array> empty_copies;
@@ -712,65 +712,6 @@ void winograd_conv_2D_gpu(
   }
 }
 
-void depthwise_conv_2D_gpu(
-    const Stream& s,
-    metal::Device& d,
-    const array& in,
-    const array& wt,
-    array out,
-    const MLXConvParams<2>& conv_params) {
-  std::ostringstream kname;
-  kname << "depthwise_conv_2d_" << type_to_name(out);
-  std::string base_name = kname.str();
-
-  const int N = conv_params.N;
-  const int ker_h = conv_params.wS[0];
-  const int ker_w = conv_params.wS[1];
-  const int str_h = conv_params.str[0];
-  const int str_w = conv_params.str[1];
-  const int tc = 8;
-  const int tw = 8;
-  const int th = 4;
-  const bool do_flip = conv_params.flip;
-
-  metal::MTLFCList func_consts = {
-      {&ker_h, MTL::DataType::DataTypeInt, 00},
-      {&ker_w, MTL::DataType::DataTypeInt, 01},
-      {&str_h, MTL::DataType::DataTypeInt, 10},
-      {&str_w, MTL::DataType::DataTypeInt, 11},
-      {&th, MTL::DataType::DataTypeInt, 100},
-      {&tw, MTL::DataType::DataTypeInt, 101},
-      {&do_flip, MTL::DataType::DataTypeBool, 200},
-  };
-
-  // clang-format off
-  kname << "_ker_h_" << ker_h
-        << "_ker_w_" << ker_w
-        << "_str_h_" << str_h
-        << "_str_w_" << str_w
-        << "_tgp_h_" << th
-        << "_tgp_w_" << tw
-        << "_do_flip_" << (do_flip ? 't' : 'n'); // clang-format on
-
-  std::string hash_name = kname.str();
-
-  auto& compute_encoder = d.get_command_encoder(s.index);
-  auto kernel = d.get_kernel(base_name, "mlx", hash_name, func_consts);
-  compute_encoder.set_compute_pipeline_state(kernel);
-
-  compute_encoder.set_input_array(in, 0);
-  compute_encoder.set_input_array(wt, 1);
-  compute_encoder.set_output_array(out, 2);
-
-  compute_encoder.set_bytes(conv_params, 3);
-
-  MTL::Size group_dims = MTL::Size(tc, tw, th);
-  MTL::Size grid_dims = MTL::Size(
-      conv_params.C / tc, conv_params.oS[1] / tw, (conv_params.oS[0] / th) * N);
-
-  compute_encoder.dispatch_threadgroups(grid_dims, group_dims);
-}
-
 void conv_2D_gpu(
     const Stream& s,
     metal::Device& d,
@@ -813,20 +754,11 @@ void conv_2D_gpu(
   bool is_kdil_one = conv_params.kdil[0] == 1 && conv_params.kdil[1] == 1;
   bool is_idil_one = conv_params.idil[0] == 1 && conv_params.idil[1] == 1;
 
-  if (is_idil_one && groups > 1) {
+  if (groups > 1) {
     const int C_per_group = conv_params.C / groups;
     const int O_per_group = conv_params.O / groups;
 
-    if (C_per_group == 1 && O_per_group == 1 && is_kdil_one &&
-        conv_params.wS[0] <= 7 && conv_params.wS[1] <= 7 &&
-        conv_params.str[0] <= 2 && conv_params.str[1] <= 2 &&
-        conv_params.oS[0] % 8 == 0 && conv_params.oS[1] % 8 == 0 &&
-        conv_params.wt_strides[1] == conv_params.wS[1] &&
-        conv_params.C % 16 == 0 && conv_params.C == conv_params.O) {
-      return depthwise_conv_2D_gpu(s, d, in, wt, out, conv_params);
-    }
-
-    if ((C_per_group <= 4 || C_per_group % 16 == 0) &&
+    if (is_idil_one && (C_per_group <= 4 || C_per_group % 16 == 0) &&
         (O_per_group <= 16 || O_per_group % 16 == 0)) {
       return implicit_gemm_conv_2D_gpu(s, d, in, wt, out, conv_params);
     } else {
@@ -923,7 +855,7 @@ void conv_3D_gpu(
 } // namespace
 
 void Convolution::eval_gpu(const std::vector<array>& inputs, array& out) {
-  out.set_data(allocator::malloc(out.nbytes()));
+  out.set_data(allocator::malloc_or_wait(out.nbytes()));
   auto& s = stream();
   auto& d = metal::device(s.device);
 

mlx/backend/metal/copy.cpp

@@ -202,7 +202,7 @@ void fill_gpu(const array& val, array& out, const Stream& s) {
   if (out.size() == 0) {
     return;
   }
-  out.set_data(allocator::malloc(out.nbytes()));
+  out.set_data(allocator::malloc_or_wait(out.nbytes()));
   bool large = out.data_size() > UINT32_MAX;
   auto& d = metal::device(s.device);
   std::string kernel_name = std::string(large ? "s2" : "s") + "_copy" +

mlx/backend/metal/custom_kernel.cpp

@@ -19,7 +19,7 @@ void CustomKernel::eval_gpu(
       copies.emplace_back(init_value_.value(), out.dtype());
       fill_gpu(copies.back(), out, s);
     } else {
-      out.set_data(allocator::malloc(out.nbytes()));
+      out.set_data(allocator::malloc_or_wait(out.nbytes()));
     }
   }
 

mlx/backend/metal/device.cpp

@@ -19,6 +19,9 @@ namespace mlx::core::metal {
 
 namespace {
 
+// TODO nicer way to set this or possibly expose as an environment variable
+constexpr int MAX_BUFFERS_PER_QUEUE = 12;
+
 constexpr const char* default_mtllib_path = METAL_PATH;
 
 auto get_metal_version() {
@@ -55,10 +58,7 @@ std::pair<MTL::Library*, NS::Error*> load_library_from_path(
 }
 
 #ifdef SWIFTPM_BUNDLE
-MTL::Library* try_load_bundle(
-    MTL::Device* device,
-    NS::URL* url,
-    const std::string& lib_name) {
+MTL::Library* try_load_bundle(MTL::Device* device, NS::URL* url) {
   std::string bundle_path = std::string(url->fileSystemRepresentation()) + "/" +
       SWIFTPM_BUNDLE + ".bundle";
   auto bundle = NS::Bundle::alloc()->init(
@@ -66,8 +66,8 @@ MTL::Library* try_load_bundle(
   if (bundle != nullptr) {
     std::string resource_path =
         std::string(bundle->resourceURL()->fileSystemRepresentation()) + "/" +
-        lib_name + ".metallib" auto [lib, error] =
-            load_library_from_path(device, resource_path.c_str());
+        "default.metallib";
+    auto [lib, error] = load_library_from_path(device, resource_path.c_str());
     if (lib) {
       return lib;
     }
@@ -76,124 +76,51 @@ MTL::Library* try_load_bundle(
 }
 #endif
 
-// Firstly, search for the metallib in the same path as this binary
-std::pair<MTL::Library*, NS::Error*> load_colocated_library(
-    MTL::Device* device,
-    const std::string& lib_name) {
-  std::string lib_path = get_colocated_mtllib_path(lib_name);
-  if (lib_path.size() != 0) {
-    return load_library_from_path(device, lib_path.c_str());
-  }
-  return {nullptr, nullptr};
-}
-
-std::pair<MTL::Library*, NS::Error*> load_swiftpm_library(
-    MTL::Device* device,
-    const std::string& lib_name) {
-#ifdef SWIFTPM_BUNDLE
-  MTL::Library* library =
-      try_load_bundle(device, NS::Bundle::mainBundle()->bundleURL(), lib_name);
-  if (library != nullptr) {
-    return {library, nullptr};
-  }
-  auto bundles = NS::Bundle::allBundles();
-  for (int i = 0, c = (int)bundles->count(); i < c; i++) {
-    auto bundle = reinterpret_cast<NS::Bundle*>(bundles->object(i));
-    library = try_load_bundle(device, bundle->resourceURL());
-    if (library != nullptr) {
-      return {library, nullptr};
-    }
-  }
-#endif
-  return {nullptr, nullptr};
-}
-
-MTL::Library* load_default_library(MTL::Device* device) {
-  NS::Error *error1, *error2, *error3;
-  MTL::Library* lib;
-  // First try the colocated mlx.metallib
-  std::tie(lib, error1) = load_colocated_library(device, "mlx");
-  if (lib) {
-    return lib;
-  }
-
-  // Then try default.metallib in a SwiftPM bundle if we have one
-  std::tie(lib, error2) = load_swiftpm_library(device, "default");
-  if (lib) {
-    return lib;
-  }
-
-  // Finally try default_mtllib_path
-  std::tie(lib, error3) = load_library_from_path(device, default_mtllib_path);
-  if (!lib) {
-    std::ostringstream msg;
-    msg << "Failed to load the default metallib. ";
-    if (error1 != nullptr) {
-      msg << error1->localizedDescription()->utf8String() << " ";
-    }
-    if (error2 != nullptr) {
-      msg << error2->localizedDescription()->utf8String() << " ";
-    }
-    if (error3 != nullptr) {
-      msg << error3->localizedDescription()->utf8String() << " ";
-    }
-    throw std::runtime_error(msg.str());
-  }
-  return lib;
-}
-
 MTL::Library* load_library(
     MTL::Device* device,
-    const std::string& lib_name,
-    const std::string& lib_path) {
-  // We have been given a path that ends in metallib so try to load it
-  if (lib_path.size() > 9 &&
-      std::equal(lib_path.end() - 9, lib_path.end(), ".metallib")) {
-    auto [lib, error] = load_library_from_path(device, lib_path.c_str());
-    if (!lib) {
-      std::ostringstream msg;
-      msg << "Failed to load the metallib from <" << lib_path << "> with error "
-          << error->localizedDescription()->utf8String();
-      throw std::runtime_error(msg.str());
-    }
-  }
-
-  // We have been given a path so try to load from lib_path / lib_name.metallib
-  if (lib_path.size() > 0) {
-    std::string full_path = lib_path + "/" + lib_name + ".metallib";
-    auto [lib, error] = load_library_from_path(device, full_path.c_str());
-    if (!lib) {
-      std::ostringstream msg;
-      msg << "Failed to load the metallib from <" << full_path
-          << "> with error " << error->localizedDescription()->utf8String();
-      throw std::runtime_error(msg.str());
-    }
-  }
-
-  // Try to load the colocated library
-  {
-    auto [lib, error] = load_colocated_library(device, lib_name);
+    const std::string& lib_name = "mlx",
+    const char* lib_path = default_mtllib_path) {
+  // Firstly, search for the metallib in the same path as this binary
+  std::string first_path = get_colocated_mtllib_path(lib_name);
+  if (first_path.size() != 0) {
+    auto [lib, error] = load_library_from_path(device, first_path.c_str());
     if (lib) {
       return lib;
     }
   }
 
-  // Try to load the library from swiftpm
-  {
-    auto [lib, error] = load_swiftpm_library(device, lib_name);
-    if (lib) {
-      return lib;
-    }
-  }
-
-  std::ostringstream msg;
-  msg << "Failed to load the metallib " << lib_name << ".metallib. "
-      << "We attempted to load it from <" << get_colocated_mtllib_path(lib_name)
-      << ">";
 #ifdef SWIFTPM_BUNDLE
-  msg << " and from the Swift PM bundle.";
+  // try to load from a swiftpm resource bundle -- scan the available bundles to
+  // find one that contains the named bundle
+  {
+    MTL::Library* library =
+        try_load_bundle(device, NS::Bundle::mainBundle()->bundleURL());
+    if (library != nullptr) {
+      return library;
+    }
+    auto bundles = NS::Bundle::allBundles();
+    for (int i = 0, c = (int)bundles->count(); i < c; i++) {
+      auto bundle = reinterpret_cast<NS::Bundle*>(bundles->object(i));
+      library = try_load_bundle(device, bundle->resourceURL());
+      if (library != nullptr) {
+        return library;
+      }
+    }
+  }
 #endif
-  throw std::runtime_error(msg.str());
+
+  // Couldn't find it so let's load it from default_mtllib_path
+  {
+    auto [lib, error] = load_library_from_path(device, lib_path);
+    if (!lib) {
+      std::ostringstream msg;
+      msg << error->localizedDescription()->utf8String() << "\n"
+          << "Failed to load device library from <" << lib_path << ">"
+          << " or <" << first_path << ">.";
+      throw std::runtime_error(msg.str());
+    }
+    return lib;
+  }
 }
 
 } // namespace
@@ -286,7 +213,7 @@ void CommandEncoder::barrier() {
 Device::Device() {
   auto pool = new_scoped_memory_pool();
   device_ = load_device();
-  library_map_ = {{"mlx", load_default_library(device_)}};
+  library_map_ = {{"mlx", load_library(device_)}};
   arch_ = std::string(device_->architecture()->name()->utf8String());
   auto arch = arch_.back();
   switch (arch) {
@@ -329,7 +256,7 @@ Device::~Device() {
 
 void Device::new_queue(int index) {
   auto thread_pool = metal::new_scoped_memory_pool();
-  auto q = device_->newCommandQueue();
+  auto q = device_->newCommandQueue(MAX_BUFFERS_PER_QUEUE);
   debug_set_stream_queue_label(q, index);
   if (!q) {
     throw std::runtime_error(

mlx/backend/metal/device.h

@@ -189,7 +189,15 @@ class Device {
 
   void register_library(
       const std::string& lib_name,
-      const std::string& lib_path = "");
+      const std::string& lib_path);
+
+  // Note, this should remain in the header so that it is not dynamically
+  // linked
+  void register_library(const std::string& lib_name) {
+    if (auto it = library_map_.find(lib_name); it == library_map_.end()) {
+      register_library(lib_name, get_colocated_mtllib_path(lib_name));
+    }
+  }
 
   MTL::Library* get_library(
       const std::string& name,

mlx/backend/metal/event.cpp

@@ -24,6 +24,10 @@ void Event::wait() {
   }
 }
 
+void Event::signal() {
+  static_cast<MTL::SharedEvent*>(event_.get())->setSignaledValue(value());
+}
+
 void Event::wait(Stream stream) {
   if (stream.device == Device::cpu) {
     scheduler::enqueue(stream, [*this]() mutable { wait(); });
@@ -38,9 +42,7 @@ void Event::wait(Stream stream) {
 
 void Event::signal(Stream stream) {
   if (stream.device == Device::cpu) {
-    scheduler::enqueue(stream, [*this]() mutable {
-      static_cast<MTL::SharedEvent*>(event_.get())->setSignaledValue(value());
-    });
+    scheduler::enqueue(stream, [*this]() mutable { signal(); });
   } else {
     auto& d = metal::device(stream.device);
     d.end_encoding(stream.index);

mlx/backend/metal/fence.cpp

@@ -20,7 +20,7 @@ struct FenceImpl {
       auto p = metal::new_scoped_memory_pool();
       fence = static_cast<void*>(d->newSharedEvent());
     } else {
-      auto buf = allocator::malloc(sizeof(uint32_t)).ptr();
+      auto buf = allocator::malloc_or_wait(sizeof(uint32_t)).ptr();
       fence = static_cast<void*>(buf);
       cpu_value()[0] = 0;
     }

mlx/backend/metal/fft.cpp

@@ -281,7 +281,7 @@ std::tuple<array, array, array> compute_raders_constants(
   }
 
   array b_q_fft({rader_n - 1}, complex64, nullptr, {});
-  b_q_fft.set_data(allocator::malloc(b_q_fft.nbytes()));
+  b_q_fft.set_data(allocator::malloc_or_wait(b_q_fft.nbytes()));
   auto b_q_fft_ptr =
       reinterpret_cast<std::complex<float>*>(b_q_fft.data<complex64_t>());
   std::ptrdiff_t item_size = b_q_fft.itemsize();
@@ -327,11 +327,11 @@ std::pair<array, array> compute_bluestein_constants(int n, int bluestein_n) {
   }
 
   array w_k({n}, complex64, nullptr, {});
-  w_k.set_data(allocator::malloc(w_k.nbytes()));
+  w_k.set_data(allocator::malloc_or_wait(w_k.nbytes()));
   std::copy(w_k_vec.begin(), w_k_vec.end(), w_k.data<complex64_t>());
 
   array w_q({bluestein_n}, complex64, nullptr, {});
-  w_q.set_data(allocator::malloc(w_q.nbytes()));
+  w_q.set_data(allocator::malloc_or_wait(w_q.nbytes()));
   auto w_q_ptr =
       reinterpret_cast<std::complex<float>*>(w_q.data<complex64_t>());
 
@@ -356,14 +356,20 @@ void multi_upload_bluestein_fft(
     bool inverse,
     bool real,
     FFTPlan& plan,
-    std::vector<array>& copies,
+    std::vector<array> copies,
     const Stream& s) {
   // TODO(alexbarron) Implement fused kernels for mutli upload bluestein's
   // algorithm
   int n = inverse ? out.shape(axis) : in.shape(axis);
   auto [w_k, w_q] = compute_bluestein_constants(n, plan.bluestein_n);
-  copies.push_back(w_k);
-  copies.push_back(w_q);
+
+  // Broadcast w_q and w_k to the batch size
+  Strides b_strides(in.ndim(), 0);
+  b_strides[axis] = 1;
+  array w_k_broadcast({}, complex64, nullptr, {});
+  array w_q_broadcast({}, complex64, nullptr, {});
+  w_k_broadcast.copy_shared_buffer(w_k, b_strides, {}, w_k.data_size());
+  w_q_broadcast.copy_shared_buffer(w_q, b_strides, {}, w_q.data_size());
 
   auto temp_shape = inverse ? out.shape() : in.shape();
   array temp(temp_shape, complex64, nullptr, {});
@@ -372,13 +378,13 @@ void multi_upload_bluestein_fft(
   if (real && !inverse) {
     // Convert float32->complex64
     copy_gpu(in, temp, CopyType::General, s);
-    copies.push_back(temp);
   } else if (real && inverse) {
     int back_offset = n % 2 == 0 ? 2 : 1;
     auto slice_shape = in.shape();
     slice_shape[axis] -= back_offset;
     array slice_temp(slice_shape, complex64, nullptr, {});
     array conj_temp(in.shape(), complex64, nullptr, {});
+    copies.push_back(slice_temp);
     copies.push_back(conj_temp);
 
     Shape rstarts(in.ndim(), 0);
@@ -388,28 +394,19 @@ void multi_upload_bluestein_fft(
     unary_op_gpu({in}, conj_temp, "Conjugate", s);
     slice_gpu(in, slice_temp, rstarts, rstrides, s);
     concatenate_gpu({conj_temp, slice_temp}, temp, (int)axis, s);
-    copies.push_back(temp);
   } else if (inverse) {
     unary_op_gpu({in}, temp, "Conjugate", s);
-    copies.push_back(temp);
   } else {
     temp.copy_shared_buffer(in);
   }
 
-  Strides b_strides(in.ndim(), 0);
-  b_strides[axis] = 1;
-  array w_k_broadcast(temp.shape(), complex64, nullptr, {});
-  w_k_broadcast.copy_shared_buffer(w_k, b_strides, {}, w_k.data_size());
   binary_op_gpu({temp, w_k_broadcast}, temp1, "Multiply", s);
 
   std::vector<std::pair<int, int>> pads;
   auto padded_shape = out.shape();
   padded_shape[axis] = plan.bluestein_n;
   array pad_temp(padded_shape, complex64, nullptr, {});
-  auto zero = array(complex64_t{0.0f, 0.0f});
-  copies.push_back(zero);
-  pad_gpu(temp1, zero, pad_temp, {(int)axis}, {0}, s);
-  copies.push_back(pad_temp);
+  pad_gpu(temp1, array(complex64_t{0.0f, 0.0f}), pad_temp, {(int)axis}, {0}, s);
 
   array pad_temp1(padded_shape, complex64, nullptr, {});
   fft_op(
@@ -421,10 +418,7 @@ void multi_upload_bluestein_fft(
       FourStepParams(),
       /*inplace=*/false,
       s);
-  copies.push_back(pad_temp1);
 
-  array w_q_broadcast(pad_temp1.shape(), complex64, nullptr, {});
-  w_q_broadcast.copy_shared_buffer(w_q, b_strides, {}, w_q.data_size());
   binary_op_gpu_inplace({pad_temp1, w_q_broadcast}, pad_temp, "Multiply", s);
 
   fft_op(
@@ -441,11 +435,9 @@ void multi_upload_bluestein_fft(
   Shape starts(in.ndim(), 0);
   Shape strides(in.ndim(), 1);
   starts[axis] = plan.bluestein_n - offset - n;
+  slice_gpu(pad_temp1, temp, starts, strides, s);
 
-  array temp2(temp_shape, complex64, nullptr, {});
-  slice_gpu(pad_temp1, temp2, starts, strides, s);
-
-  binary_op_gpu_inplace({temp2, w_k_broadcast}, temp1, "Multiply", s);
+  binary_op_gpu_inplace({temp, w_k_broadcast}, temp1, "Multiply", s);
 
   if (real && !inverse) {
     Shape rstarts(in.ndim(), 0);
@@ -457,21 +449,26 @@ void multi_upload_bluestein_fft(
     array temp_float(out.shape(), out.dtype(), nullptr, {});
     copies.push_back(temp_float);
     copies.push_back(inv_n);
-    copies.push_back(temp1);
 
     copy_gpu(temp1, temp_float, CopyType::General, s);
     binary_op_gpu({temp_float, inv_n}, out, "Multiply", s);
   } else if (inverse) {
     auto inv_n = array({1.0f / n}, {1}, complex64);
-    array temp3(temp_shape, complex64, nullptr, {});
-    unary_op_gpu({temp1}, temp3, "Conjugate", s);
-    binary_op_gpu({temp3, inv_n}, out, "Multiply", s);
+    unary_op_gpu({temp1}, temp, "Conjugate", s);
+    binary_op_gpu({temp, inv_n}, out, "Multiply", s);
     copies.push_back(inv_n);
-    copies.push_back(temp1);
-    copies.push_back(temp3);
   } else {
     out.copy_shared_buffer(temp1);
   }
+
+  copies.push_back(w_k);
+  copies.push_back(w_q);
+  copies.push_back(w_k_broadcast);
+  copies.push_back(w_q_broadcast);
+  copies.push_back(temp);
+  copies.push_back(temp1);
+  copies.push_back(pad_temp);
+  copies.push_back(pad_temp1);
 }
 
 void four_step_fft(
@@ -481,9 +478,8 @@ void four_step_fft(
     bool inverse,
     bool real,
     FFTPlan& plan,
-    std::vector<array>& copies,
-    const Stream& s,
-    bool in_place) {
+    std::vector<array> copies,
+    const Stream& s) {
   auto& d = metal::device(s.device);
 
   if (plan.bluestein_n == -1) {
@@ -496,14 +492,7 @@ void four_step_fft(
         in, temp, axis, inverse, real, four_step_params, /*inplace=*/false, s);
     four_step_params.first_step = false;
     fft_op(
-        temp,
-        out,
-        axis,
-        inverse,
-        real,
-        four_step_params,
-        /*inplace=*/in_place,
-        s);
+        temp, out, axis, inverse, real, four_step_params, /*inplace=*/false, s);
     copies.push_back(temp);
   } else {
     multi_upload_bluestein_fft(in, out, axis, inverse, real, plan, copies, s);
@@ -562,7 +551,8 @@ void fft_op(
       flags.row_contiguous = is_row_contiguous;
       flags.contiguous = data_size == x_copy.size();
 
-      x_copy.set_data(allocator::malloc(x.nbytes()), data_size, strides, flags);
+      x_copy.set_data(
+          allocator::malloc_or_wait(x.nbytes()), data_size, strides, flags);
       copy_gpu_inplace(x, x_copy, CopyType::GeneralGeneral, s);
       copies.push_back(x_copy);
       return x_copy;
@@ -585,7 +575,7 @@ void fft_op(
 
   auto plan = plan_fft(n);
   if (plan.four_step) {
-    four_step_fft(in, out, axis, inverse, real, plan, copies, s, inplace);
+    four_step_fft(in, out, axis, inverse, real, plan, copies, s);
     d.add_temporaries(std::move(copies), s.index);
     return;
   }
@@ -593,7 +583,7 @@ void fft_op(
   // TODO: allow donation here
   if (!inplace) {
     out.set_data(
-        allocator::malloc(out.nbytes()),
+        allocator::malloc_or_wait(out.nbytes()),
         out_data_size,
         out_strides,
         in_contiguous.flags());

mlx/backend/metal/hadamard.cpp

@@ -84,7 +84,7 @@ void Hadamard::eval_gpu(const std::vector<array>& inputs, array& out) {
   if (in_contiguous.is_donatable()) {
     out.copy_shared_buffer(in_contiguous);
   } else {
-    out.set_data(allocator::malloc(out.nbytes()));
+    out.set_data(allocator::malloc_or_wait(out.nbytes()));
   }
 
   int n, m;
@@ -161,7 +161,7 @@ void Hadamard::eval_gpu(const std::vector<array>& inputs, array& out) {
     // Upload 2:
     // y = h12 @ tmp
     array temp(in.shape(), in.dtype(), nullptr, {});
-    temp.set_data(allocator::malloc(temp.nbytes()));
+    temp.set_data(allocator::malloc_or_wait(temp.nbytes()));
     copies.push_back(temp);
 
     launch_hadamard(in_contiguous, temp, "n" + kernel_name, 1.0);

mlx/backend/metal/indexing.cpp

@@ -43,7 +43,7 @@ void Gather::eval_gpu(const std::vector<array>& inputs, array& out) {
     throw std::runtime_error(msg.str());
   }
 
-  out.set_data(allocator::malloc(out.nbytes()));
+  out.set_data(allocator::malloc_or_wait(out.nbytes()));
   if (out.size() == 0) {
     return;
   }
@@ -393,7 +393,7 @@ void GatherAxis::eval_gpu(const std::vector<array>& inputs, array& out) {
   auto& src = inputs[0];
   auto& idx = inputs[1];
 
-  out.set_data(allocator::malloc(out.nbytes()));
+  out.set_data(allocator::malloc_or_wait(out.nbytes()));
   if (out.size() == 0) {
     return;
   }

mlx/backend/metal/jit/arange.h

@@ -0,0 +1,9 @@
+// Copyright © 2024 Apple Inc.
+
+constexpr std::string_view arange_kernels = R"(
+template [[host_name("{0}")]] [[kernel]] void arange<{1}>(
+    constant const {1}& start,
+    constant const {1}& step,
+    device {1}* out,
+    uint index [[thread_position_in_grid]]);
+)";

mlx/backend/metal/jit/includes.h

@@ -20,7 +20,6 @@ const char* copy();
 const char* fft();
 const char* gather_axis();
 const char* hadamard();
-const char* logsumexp();
 const char* quantized();
 const char* ternary();
 const char* scan();
@@ -33,7 +32,6 @@ const char* gemm();
 const char* steel_gemm_fused();
 const char* steel_gemm_masked();
 const char* steel_gemm_splitk();
-const char* steel_gemm_gather();
 const char* conv();
 const char* steel_conv();
 const char* steel_conv_general();

mlx/backend/metal/jit/softmax.h

@@ -0,0 +1,23 @@
+// Copyright © 2024 Apple Inc.
+
+constexpr std::string_view softmax_kernels = R"(
+template [[host_name("block_{0}")]] [[kernel]] void
+softmax_single_row<{1}, {2}>(
+    const device {1}* in,
+    device {1}* out,
+    constant int& axis_size,
+    uint gid [[thread_position_in_grid]],
+    uint _lid [[thread_position_in_threadgroup]],
+    uint simd_lane_id [[thread_index_in_simdgroup]],
+    uint simd_group_id [[simdgroup_index_in_threadgroup]]);
+template [[host_name("looped_{0}")]] [[kernel]] void
+softmax_looped<{1}, {2}>(
+    const device {1}* in,
+    device {1}* out,
+    constant int& axis_size,
+    uint gid [[threadgroup_position_in_grid]],
+    uint lid [[thread_position_in_threadgroup]],
+    uint lsize [[threads_per_threadgroup]],
+    uint simd_lane_id [[thread_index_in_simdgroup]],
+    uint simd_group_id [[simdgroup_index_in_threadgroup]]);
+)";

mlx/backend/metal/jit_kernels.cpp

@@ -1,6 +1,8 @@
 // Copyright © 2024 Apple Inc.
 #include "mlx/backend/common/compiled.h"
+#include "mlx/backend/metal/jit/arange.h"
 #include "mlx/backend/metal/jit/includes.h"
+#include "mlx/backend/metal/jit/softmax.h"
 #include "mlx/backend/metal/kernels.h"
 #include "mlx/backend/metal/utils.h"
 
@@ -19,11 +21,13 @@ MTL::ComputePipelineState* get_arange_kernel(
     const std::string& kernel_name,
     const array& out) {
   auto lib = d.get_library(kernel_name, [&]() {
-    std::string kernel_source = metal::utils();
-    kernel_source += metal::arange();
-    kernel_source += get_template_definition(
-        kernel_name, "arange", get_type_string(out.dtype()));
-    return kernel_source;
+    std::ostringstream kernel_source;
+    kernel_source << metal::utils() << metal::arange()
+                  << fmt::format(
+                         arange_kernels,
+                         kernel_name,
+                         get_type_string(out.dtype()));
+    return kernel_source.str();
   });
   return d.get_kernel(kernel_name, lib);
 }
@@ -255,34 +259,14 @@ MTL::ComputePipelineState* get_softmax_kernel(
     const array& out) {
   std::string lib_name = kernel_name.substr(kernel_name.find("_") + 1);
   auto lib = d.get_library(lib_name, [&] {
-    std::string kernel_source = metal::utils();
-    auto in_type = get_type_string(out.dtype());
-    auto acc_type = get_type_string(precise ? float32 : out.dtype());
-    kernel_source += metal::softmax();
-    kernel_source += get_template_definition(
-        "block_" + lib_name, "softmax_single_row", in_type, acc_type);
-    kernel_source += get_template_definition(
-        "looped_" + lib_name, "softmax_looped", in_type, acc_type);
-    return kernel_source;
-  });
-  return d.get_kernel(kernel_name, lib);
-}
-
-MTL::ComputePipelineState* get_logsumexp_kernel(
-    metal::Device& d,
-    const std::string& kernel_name,
-    const array& out) {
-  std::string lib_name = kernel_name.substr(kernel_name.find("_") + 1);
-  auto lib = d.get_library(lib_name, [&] {
-    auto t_str = get_type_string(out.dtype());
-    std::string kernel_source;
-    kernel_source = metal::utils();
-    kernel_source += metal::logsumexp();
-    kernel_source +=
-        get_template_definition("block_" + lib_name, "logsumexp", t_str);
-    kernel_source += get_template_definition(
-        "looped_" + lib_name, "logsumexp_looped", t_str);
-    return kernel_source;
+    std::ostringstream kernel_source;
+    kernel_source << metal::utils() << metal::softmax()
+                  << fmt::format(
+                         softmax_kernels,
+                         lib_name,
+                         get_type_string(out.dtype()),
+                         get_type_string(precise ? float32 : out.dtype()));
+    return kernel_source.str();
   });
   return d.get_kernel(kernel_name, lib);
 }
@@ -584,44 +568,6 @@ MTL::ComputePipelineState* get_steel_gemm_masked_kernel(
   return d.get_kernel(kernel_name, lib);
 }
 
-MTL::ComputePipelineState* get_steel_gemm_gather_kernel(
-    metal::Device& d,
-    const std::string& kernel_name,
-    const std::string& hash_name,
-    const metal::MTLFCList& func_consts,
-    const array& out,
-    bool transpose_a,
-    bool transpose_b,
-    int bm,
-    int bn,
-    int bk,
-    int wm,
-    int wn,
-    bool rhs) {
-  const auto& lib_name = kernel_name;
-  auto lib = d.get_library(lib_name, [&]() {
-    std::string kernel_source;
-    concatenate(
-        kernel_source,
-        metal::utils(),
-        metal::gemm(),
-        metal::steel_gemm_gather(),
-        get_template_definition(
-            lib_name,
-            rhs ? "gather_mm_rhs" : "gather_mm",
-            get_type_string(out.dtype()),
-            bm,
-            bn,
-            bk,
-            wm,
-            wn,
-            transpose_a,
-            transpose_b));
-    return kernel_source;
-  });
-  return d.get_kernel(kernel_name, lib, hash_name, func_consts);
-}
-
 MTL::ComputePipelineState* get_gemv_masked_kernel(
     metal::Device& d,
     const std::string& kernel_name,

mlx/backend/metal/kernels.h

@@ -59,11 +59,6 @@ MTL::ComputePipelineState* get_softmax_kernel(
     bool precise,
     const array& out);
 
-MTL::ComputePipelineState* get_logsumexp_kernel(
-    metal::Device& d,
-    const std::string& kernel_name,
-    const array& out);
-
 MTL::ComputePipelineState* get_scan_kernel(
     metal::Device& d,
     const std::string& kernel_name,
@@ -160,21 +155,6 @@ MTL::ComputePipelineState* get_steel_gemm_masked_kernel(
     bool mn_aligned,
     bool k_aligned);
 
-MTL::ComputePipelineState* get_steel_gemm_gather_kernel(
-    metal::Device& d,
-    const std::string& kernel_name,
-    const std::string& hash_name,
-    const metal::MTLFCList& func_consts,
-    const array& out,
-    bool transpose_a,
-    bool transpose_b,
-    int bm,
-    int bn,
-    int bk,
-    int wm,
-    int wn,
-    bool rhs);
-
 MTL::ComputePipelineState* get_steel_conv_kernel(
     metal::Device& d,
     const std::string& kernel_name,

mlx/backend/metal/kernels/CMakeLists.txt

@@ -13,10 +13,6 @@ function(build_kernel_base TARGET SRCFILE DEPS)
   if(MLX_METAL_DEBUG)
     set(METAL_FLAGS ${METAL_FLAGS} -gline-tables-only -frecord-sources)
   endif()
-  if(NOT CMAKE_OSX_DEPLOYMENT_TARGET STREQUAL "")
-    set(METAL_FLAGS ${METAL_FLAGS}
-                    "-mmacosx-version-min=${CMAKE_OSX_DEPLOYMENT_TARGET}")
-  endif()
   if(MLX_METAL_VERSION GREATER_EQUAL 310)
     set(VERSION_INCLUDES
         ${PROJECT_SOURCE_DIR}/mlx/backend/metal/kernels/metal_3_1)
@@ -69,7 +65,6 @@ set(STEEL_HEADERS
     steel/gemm/loader.h
     steel/gemm/transforms.h
     steel/gemm/kernels/steel_gemm_fused.h
-    steel/gemm/kernels/steel_gemm_gather.h
     steel/gemm/kernels/steel_gemm_masked.h
     steel/gemm/kernels/steel_gemm_splitk.h
     steel/utils/type_traits.h
@@ -110,14 +105,12 @@ if(NOT MLX_METAL_JIT)
   build_kernel(quantized quantized.h ${STEEL_HEADERS})
   build_kernel(scan scan.h)
   build_kernel(softmax softmax.h)
-  build_kernel(logsumexp logsumexp.h)
   build_kernel(sort sort.h)
   build_kernel(ternary ternary.h ternary_ops.h)
   build_kernel(unary unary.h unary_ops.h)
   build_kernel(steel/conv/kernels/steel_conv ${STEEL_HEADERS})
   build_kernel(steel/conv/kernels/steel_conv_general ${STEEL_HEADERS})
   build_kernel(steel/gemm/kernels/steel_gemm_fused ${STEEL_HEADERS})
-  build_kernel(steel/gemm/kernels/steel_gemm_gather ${STEEL_HEADERS})
   build_kernel(steel/gemm/kernels/steel_gemm_masked ${STEEL_HEADERS})
   build_kernel(steel/gemm/kernels/steel_gemm_splitk ${STEEL_HEADERS})
   build_kernel(gemv_masked steel/utils.h)

mlx/backend/metal/kernels/arange.metal

@@ -5,7 +5,11 @@
 #include "mlx/backend/metal/kernels/arange.h"
 
 #define instantiate_arange(tname, type)                                 \
-  instantiate_kernel("arange" #tname, arange, type)
+  template [[host_name("arange" #tname)]] [[kernel]] void arange<type>( \
+      constant const type& start,                                       \
+      constant const type& step,                                        \
+      device type* out,                                                 \
+      uint index [[thread_position_in_grid]]);
 
 instantiate_arange(uint8, uint8_t)
 instantiate_arange(uint16, uint16_t)

mlx/backend/metal/kernels/conv.metal

@@ -275,128 +275,6 @@ instantiate_naive_conv_2d_blocks(float32, float);
 instantiate_naive_conv_2d_blocks(float16, half);
 instantiate_naive_conv_2d_blocks(bfloat16, bfloat16_t);
 
-///////////////////////////////////////////////////////////////////////////////
-/// Depthwise convolution kernels
-///////////////////////////////////////////////////////////////////////////////
-
-constant int ker_h [[function_constant(00)]];
-constant int ker_w [[function_constant(01)]];
-constant int str_h [[function_constant(10)]];
-constant int str_w [[function_constant(11)]];
-constant int tgp_h [[function_constant(100)]];
-constant int tgp_w [[function_constant(101)]];
-constant bool do_flip [[function_constant(200)]];
-
-constant int span_h = tgp_h * str_h + ker_h - 1;
-constant int span_w = tgp_w * str_w + ker_w - 1;
-constant int span_hw = span_h * span_w;
-
-template <typename T>
-[[kernel]] void depthwise_conv_2d(
-    const device T* in [[buffer(0)]],
-    const device T* wt [[buffer(1)]],
-    device T* out [[buffer(2)]],
-    const constant MLXConvParams<2>& params [[buffer(3)]],
-    uint3 tid [[threadgroup_position_in_grid]],
-    uint3 lid [[thread_position_in_threadgroup]],
-    uint3 gid [[thread_position_in_grid]],
-    uint simd_gid [[simdgroup_index_in_threadgroup]],
-    uint simd_lid [[thread_index_in_simdgroup]]) {
-  constexpr int tc = 8;
-  constexpr int tw = 8;
-  constexpr int th = 4;
-
-  constexpr int c_per_thr = 8;
-
-  constexpr int TGH = th * 2 + 6;
-  constexpr int TGW = tw * 2 + 6;
-  constexpr int TGC = tc;
-
-  threadgroup T ins[TGH * TGW * TGC];
-
-  const int n_tgblocks_h = params.oS[0] / th;
-  const int n = tid.z / n_tgblocks_h;
-  const int tghid = tid.z % n_tgblocks_h;
-  const int oh = tghid * th + lid.z;
-  const int ow = gid.y;
-  const int c = gid.x;
-
-  in += n * params.in_strides[0];
-
-  // Load in
-  {
-    constexpr int n_threads = th * tw * tc;
-    const int tg_oh = (tghid * th) * str_h - params.pad[0];
-    const int tg_ow = (tid.y * tw) * str_w - params.pad[1];
-    const int tg_c = tid.x * tc;
-
-    const int thread_idx = simd_gid * 32 + simd_lid;
-    constexpr int thr_per_hw = tc / c_per_thr;
-    constexpr int hw_per_group = n_threads / thr_per_hw;
-
-    const int thr_c = thread_idx % thr_per_hw;
-    const int thr_hw = thread_idx / thr_per_hw;
-
-    for (int hw = thr_hw; hw < span_hw; hw += hw_per_group) {
-      const int h = hw / span_w;
-      const int w = hw % span_w;
-
-      const int ih = tg_oh + h;
-      const int iw = tg_ow + w;
-
-      const int in_s_offset = h * span_w * TGC + w * TGC;
-
-      if (ih >= 0 && ih < params.iS[0] && iw >= 0 && iw < params.iS[1]) {
-        const auto in_load =
-            in + ih * params.in_strides[1] + iw * params.in_strides[2] + tg_c;
-
-        MLX_MTL_PRAGMA_UNROLL
-        for (int cc = 0; cc < c_per_thr; ++cc) {
-          ins[in_s_offset + c_per_thr * thr_c + cc] =
-              in_load[c_per_thr * thr_c + cc];
-        }
-      } else {
-        MLX_MTL_PRAGMA_UNROLL
-        for (int cc = 0; cc < c_per_thr; ++cc) {
-          ins[in_s_offset + c_per_thr * thr_c + cc] = T(0);
-        }
-      }
-    }
-  }
-
-  threadgroup_barrier(mem_flags::mem_threadgroup);
-  wt += c * params.wt_strides[0];
-
-  const auto ins_ptr =
-      &ins[lid.z * str_h * span_w * TGC + lid.y * str_w * TGC + lid.x];
-  float o = 0.;
-  for (int h = 0; h < ker_h; ++h) {
-    for (int w = 0; w < ker_w; ++w) {
-      int wt_h = h;
-      int wt_w = w;
-      if (do_flip) {
-        wt_h = ker_h - h - 1;
-        wt_w = ker_w - w - 1;
-      }
-      auto inv = ins_ptr[h * span_w * TGC + w * TGC];
-      auto wtv = wt[wt_h * ker_w + wt_w];
-      o += inv * wtv;
-    }
-  }
-  threadgroup_barrier(mem_flags::mem_none);
-
-  out += n * params.out_strides[0] + oh * params.out_strides[1] +
-      ow * params.out_strides[2];
-  out[c] = static_cast<T>(o);
-}
-
-#define instantiate_depthconv2d(iname, itype) \
-  instantiate_kernel("depthwise_conv_2d_" #iname, depthwise_conv_2d, itype)
-
-instantiate_depthconv2d(float32, float);
-instantiate_depthconv2d(float16, half);
-instantiate_depthconv2d(bfloat16, bfloat16_t);
-
 ///////////////////////////////////////////////////////////////////////////////
 /// Winograd kernels
 ///////////////////////////////////////////////////////////////////////////////

mlx/backend/metal/kernels/fft.h

@@ -483,4 +483,4 @@ template <
   perform_fft(fft_idx, &p, m, n, buf);
 
   read_writer.write_strided(stride, overall_n);
-}
+}

mlx/backend/metal/kernels/gemv.metal

@@ -341,7 +341,7 @@ struct GEMVTKernel {
 
         MLX_MTL_PRAGMA_UNROLL
         for (int tm = 0; tm < TM; tm++) {
-          auto vc = static_cast<AccT>(v_coeff[tm]);
+          auto vc = float(v_coeff[tm]);
           for (int tn = 0; tn < TN; tn++) {
             inter[tn] = mat[(bm + tm) * marix_ld + out_col + tn];
           }

mlx/backend/metal/kernels/layer_norm.metal

@@ -493,11 +493,71 @@ template <typename T, int N_READS = RMS_N_READS>
 }
 
 // clang-format off
-#define instantiate_layer_norm(name, itype)                                       \
-  instantiate_kernel("layer_norm" #name, layer_norm_single_row, itype)            \
-  instantiate_kernel("vjp_layer_norm" #name, vjp_layer_norm_single_row, itype)    \
-  instantiate_kernel("layer_norm_looped" #name, layer_norm_looped, itype)         \
-  instantiate_kernel("vjp_layer_norm_looped" #name, vjp_layer_norm_looped, itype)
+#define instantiate_layer_norm_single_row(name, itype)            \
+  template [[host_name("layer_norm" #name)]] [[kernel]] void      \
+  layer_norm_single_row<itype>(                                   \
+      const device itype* x,                                      \
+      const device itype* w,                                      \
+      const device itype* b,                                      \
+      device itype* out,                                          \
+      constant float& eps,                                        \
+      constant uint& axis_size,                                   \
+      constant uint& w_stride,                                    \
+      constant uint& b_stride,                                    \
+      uint gid [[thread_position_in_grid]],                       \
+      uint lid [[thread_position_in_threadgroup]],                \
+      uint simd_lane_id [[thread_index_in_simdgroup]],            \
+      uint simd_group_id [[simdgroup_index_in_threadgroup]]);     \
+  template [[host_name("vjp_layer_norm" #name)]] [[kernel]] void  \
+  vjp_layer_norm_single_row<itype>(                               \
+      const device itype* x,                                      \
+      const device itype* w,                                      \
+      const device itype* g,                                      \
+      device itype* gx,                                           \
+      device itype* gw,                                           \
+      constant float& eps,                                        \
+      constant uint& axis_size,                                   \
+      constant uint& w_stride,                                    \
+      uint gid [[thread_position_in_grid]],                       \
+      uint lid [[thread_position_in_threadgroup]],                \
+      uint simd_lane_id [[thread_index_in_simdgroup]],            \
+      uint simd_group_id [[simdgroup_index_in_threadgroup]]);
+
+#define instantiate_layer_norm_looped(name, itype)                       \
+  template [[host_name("layer_norm_looped" #name)]] [[kernel]] void      \
+  layer_norm_looped<itype>(                                              \
+      const device itype* x,                                             \
+      const device itype* w,                                             \
+      const device itype* b,                                             \
+      device itype* out,                                                 \
+      constant float& eps,                                               \
+      constant uint& axis_size,                                          \
+      constant uint& w_stride,                                           \
+      constant uint& b_stride,                                           \
+      uint gid [[thread_position_in_grid]],                              \
+      uint lid [[thread_position_in_threadgroup]],                       \
+      uint lsize [[threads_per_threadgroup]],                            \
+      uint simd_lane_id [[thread_index_in_simdgroup]],                   \
+      uint simd_group_id [[simdgroup_index_in_threadgroup]]);            \
+  template [[host_name("vjp_layer_norm_looped" #name)]] [[kernel]] void  \
+  vjp_layer_norm_looped<itype>(                                          \
+      const device itype* x,                                             \
+      const device itype* w,                                             \
+      const device itype* g,                                             \
+      device itype* gx,                                                  \
+      device itype* gb,                                                  \
+      constant float& eps,                                               \
+      constant uint& axis_size,                                          \
+      constant uint& w_stride,                                           \
+      uint gid [[thread_position_in_grid]],                              \
+      uint lid [[thread_position_in_threadgroup]],                       \
+      uint lsize [[threads_per_threadgroup]],                            \
+      uint simd_lane_id [[thread_index_in_simdgroup]],                   \
+      uint simd_group_id [[simdgroup_index_in_threadgroup]]);
+
+#define instantiate_layer_norm(name, itype)      \
+  instantiate_layer_norm_single_row(name, itype) \
+  instantiate_layer_norm_looped(name, itype)
 
 instantiate_layer_norm(float32, float)
 instantiate_layer_norm(float16, half)

mlx/backend/metal/kernels/logsumexp.h

@@ -1,143 +0,0 @@
-// Copyright © 2025 Apple Inc.
-
-template <typename T, typename AccT = float, int N_READS = 4>
-[[kernel]] void logsumexp(
-    const device T* in,
-    device T* out,
-    constant int& axis_size,
-    uint gid [[threadgroup_position_in_grid]],
-    uint _lid [[thread_position_in_threadgroup]],
-    uint simd_lane_id [[thread_index_in_simdgroup]],
-    uint simd_group_id [[simdgroup_index_in_threadgroup]]) {
-  int lid = _lid;
-
-  constexpr int SIMD_SIZE = 32;
-
-  threadgroup AccT local_max[SIMD_SIZE];
-  threadgroup AccT local_normalizer[SIMD_SIZE];
-
-  AccT ld[N_READS];
-
-  in += gid * size_t(axis_size) + lid * N_READS;
-  if (lid * N_READS + N_READS <= axis_size) {
-    for (int i = 0; i < N_READS; i++) {
-      ld[i] = AccT(in[i]);
-    }
-  } else {
-    for (int i = 0; i < N_READS; i++) {
-      ld[i] =
-          ((lid * N_READS + i) < axis_size) ? AccT(in[i]) : Limits<AccT>::min;
-    }
-  }
-  if (simd_group_id == 0) {
-    local_max[simd_lane_id] = Limits<AccT>::min;
-    local_normalizer[simd_lane_id] = 0;
-  }
-  threadgroup_barrier(mem_flags::mem_threadgroup);
-
-  // Get the max
-  AccT maxval = Limits<AccT>::finite_min;
-  for (int i = 0; i < N_READS; i++) {
-    maxval = (maxval < ld[i]) ? ld[i] : maxval;
-  }
-  maxval = simd_max(maxval);
-  if (simd_lane_id == 0) {
-    local_max[simd_group_id] = maxval;
-  }
-  threadgroup_barrier(mem_flags::mem_threadgroup);
-  if (simd_group_id == 0) {
-    maxval = simd_max(local_max[simd_lane_id]);
-    if (simd_lane_id == 0) {
-      local_max[0] = maxval;
-    }
-  }
-  threadgroup_barrier(mem_flags::mem_threadgroup);
-  maxval = local_max[0];
-
-  // Compute exp(x_i - maxval) and store the partial sums in local_normalizer
-  AccT normalizer = 0;
-  for (int i = 0; i < N_READS; i++) {
-    normalizer += fast::exp(ld[i] - maxval);
-  }
-  normalizer = simd_sum(normalizer);
-  if (simd_lane_id == 0) {
-    local_normalizer[simd_group_id] = normalizer;
-  }
-  threadgroup_barrier(mem_flags::mem_threadgroup);
-  if (simd_group_id == 0) {
-    normalizer = simd_sum(local_normalizer[simd_lane_id]);
-    if (simd_lane_id == 0) {
-      out[gid] = isinf(maxval) ? T(maxval) : T(log(normalizer) + maxval);
-    }
-  }
-}
-
-template <typename T, typename AccT = float, int N_READS = 4>
-[[kernel]] void logsumexp_looped(
-    const device T* in,
-    device T* out,
-    constant int& axis_size,
-    uint gid [[threadgroup_position_in_grid]],
-    uint lid [[thread_position_in_threadgroup]],
-    uint lsize [[threads_per_threadgroup]],
-    uint simd_lane_id [[thread_index_in_simdgroup]],
-    uint simd_group_id [[simdgroup_index_in_threadgroup]]) {
-  in += gid * size_t(axis_size);
-
-  constexpr int SIMD_SIZE = 32;
-
-  threadgroup AccT local_max[SIMD_SIZE];
-  threadgroup AccT local_normalizer[SIMD_SIZE];
-
-  // Get the max and the normalizer in one go
-  AccT prevmax;
-  AccT maxval = Limits<AccT>::finite_min;
-  AccT normalizer = 0;
-  for (int r = 0; r < static_cast<int>(ceildiv(axis_size, N_READS * lsize));
-       r++) {
-    int offset = r * lsize * N_READS + lid * N_READS;
-    AccT vals[N_READS];
-    if (offset + N_READS <= axis_size) {
-      for (int i = 0; i < N_READS; i++) {
-        vals[i] = AccT(in[offset + i]);
-      }
-    } else {
-      for (int i = 0; i < N_READS; i++) {
-        vals[i] = (offset + i < axis_size) ? AccT(in[offset + i])
-                                           : Limits<AccT>::finite_min;
-      }
-    }
-    prevmax = maxval;
-    for (int i = 0; i < N_READS; i++) {
-      maxval = (maxval < vals[i]) ? vals[i] : maxval;
-    }
-    normalizer *= fast::exp(prevmax - maxval);
-    for (int i = 0; i < N_READS; i++) {
-      normalizer += fast::exp(vals[i] - maxval);
-    }
-  }
-  prevmax = maxval;
-  maxval = simd_max(maxval);
-  normalizer *= fast::exp(prevmax - maxval);
-  normalizer = simd_sum(normalizer);
-
-  prevmax = maxval;
-  if (simd_lane_id == 0) {
-    local_max[simd_group_id] = maxval;
-  }
-  threadgroup_barrier(mem_flags::mem_threadgroup);
-  maxval = simd_max(local_max[simd_lane_id]);
-  normalizer *= fast::exp(prevmax - maxval);
-  if (simd_lane_id == 0) {
-    local_normalizer[simd_group_id] = normalizer;
-  }
-  threadgroup_barrier(mem_flags::mem_threadgroup);
-  normalizer = simd_sum(local_normalizer[simd_lane_id]);
-
-  if (simd_group_id == 0) {
-    normalizer = simd_sum(local_normalizer[simd_lane_id]);
-    if (simd_lane_id == 0) {
-      out[gid] = isinf(maxval) ? T(maxval) : T(log(normalizer) + maxval);
-    }
-  }
-}

mlx/backend/metal/kernels/logsumexp.metal

@@ -1,18 +0,0 @@
-// Copyright © 2023-2024 Apple Inc.
-
-#include <metal_common>
-#include <metal_simdgroup>
-
-using namespace metal;
-
-// clang-format off
-#include "mlx/backend/metal/kernels/utils.h"
-#include "mlx/backend/metal/kernels/logsumexp.h"
-
-#define instantiate_logsumexp(name, itype)                               \
-  instantiate_kernel("block_logsumexp_" #name, logsumexp, itype)         \
-  instantiate_kernel("looped_logsumexp_" #name, logsumexp_looped, itype) \
-
-instantiate_logsumexp(float32, float)
-instantiate_logsumexp(float16, half)
-instantiate_logsumexp(bfloat16, bfloat16_t) // clang-format on

mlx/backend/metal/kernels/quantized.h

@@ -586,13 +586,13 @@ METAL_FUNC void qmv_quad_impl(
   // Adjust positions
   const int in_vec_size_w = in_vec_size / pack_factor;
   const int in_vec_size_g = in_vec_size / group_size;
-  const int out_row = tid.y * quads_per_simd * results_per_quadgroup + quad_gid;
+  const int out_row = tid.x * quads_per_simd * results_per_quadgroup + quad_gid;
 
   w += out_row * in_vec_size_w + quad_lid * packs_per_thread;
   scales += out_row * in_vec_size_g + quad_lid / scale_step_per_thread;
   biases += out_row * in_vec_size_g + quad_lid / scale_step_per_thread;
-  x += tid.x * in_vec_size + quad_lid * values_per_thread;
-  y += tid.x * out_vec_size + out_row;
+  x += tid.y * in_vec_size + quad_lid * values_per_thread;
+  y += tid.y * out_vec_size + out_row;
 
   U sum = load_vector<T, U, values_per_thread, bits>(x, x_thread);
 

mlx/backend/metal/kernels/rms_norm.metal

@@ -380,11 +380,69 @@ template <typename T, int N_READS = RMS_N_READS>
 }
 
 // clang-format off
-#define instantiate_rms(name, itype)                                \
-  instantiate_kernel("rms" #name, rms_single_row, itype)            \
-  instantiate_kernel("vjp_rms" #name, vjp_rms_single_row, itype)    \
-  instantiate_kernel("rms_looped" #name, rms_looped, itype)         \
-  instantiate_kernel("vjp_rms_looped" #name, vjp_rms_looped, itype)
+#define instantiate_rms_single_row(name, itype)               \
+  template [[host_name("rms" #name)]] [[kernel]] void         \
+  rms_single_row<itype>(                                      \
+      const device itype* x,                                  \
+      const device itype* w,                                  \
+      device itype* out,                                      \
+      constant float& eps,                                    \
+      constant uint& axis_size,                               \
+      constant uint& w_stride,                                \
+      uint gid [[thread_position_in_grid]],                   \
+      uint lid [[thread_position_in_threadgroup]],            \
+      uint simd_lane_id [[thread_index_in_simdgroup]],        \
+      uint simd_group_id [[simdgroup_index_in_threadgroup]]); \
+                                                              \
+  template [[host_name("vjp_rms" #name)]] [[kernel]] void     \
+  vjp_rms_single_row<itype>(                                  \
+      const device itype* x,                                  \
+      const device itype* w,                                  \
+      const device itype* g,                                  \
+      device itype* gx,                                       \
+      device itype* gw,                                       \
+      constant float& eps,                                    \
+      constant uint& axis_size,                               \
+      constant uint& w_stride,                                \
+      uint gid [[thread_position_in_grid]],                   \
+      uint lid [[thread_position_in_threadgroup]],            \
+      uint simd_lane_id [[thread_index_in_simdgroup]],        \
+      uint simd_group_id [[simdgroup_index_in_threadgroup]]);
+
+#define instantiate_rms_looped(name, itype)                      \
+  template [[host_name("rms_looped" #name)]] [[kernel]] void     \
+  rms_looped<itype>(                                             \
+      const device itype* x,                                     \
+      const device itype* w,                                     \
+      device itype* out,                                         \
+      constant float& eps,                                       \
+      constant uint& axis_size,                                  \
+      constant uint& w_stride,                                   \
+      uint gid [[thread_position_in_grid]],                      \
+      uint lid [[thread_position_in_threadgroup]],               \
+      uint lsize [[threads_per_threadgroup]],                    \
+      uint simd_lane_id [[thread_index_in_simdgroup]],           \
+      uint simd_group_id [[simdgroup_index_in_threadgroup]]);    \
+                                                                 \
+  template [[host_name("vjp_rms_looped" #name)]] [[kernel]] void \
+  vjp_rms_looped<itype>(                                         \
+      const device itype* x,                                     \
+      const device itype* w,                                     \
+      const device itype* g,                                     \
+      device itype* gx,                                          \
+      device itype* gw,                                          \
+      constant float& eps,                                       \
+      constant uint& axis_size,                                  \
+      constant uint& w_stride,                                   \
+      uint gid [[thread_position_in_grid]],                      \
+      uint lid [[thread_position_in_threadgroup]],               \
+      uint lsize [[threads_per_threadgroup]],                    \
+      uint simd_lane_id [[thread_index_in_simdgroup]],           \
+      uint simd_group_id [[simdgroup_index_in_threadgroup]]);
+
+#define instantiate_rms(name, itype)      \
+  instantiate_rms_single_row(name, itype) \
+  instantiate_rms_looped(name, itype)
 
 instantiate_rms(float32, float)
 instantiate_rms(float16, half)

mlx/backend/metal/kernels/scaled_dot_product_attention.metal

@@ -1,11 +1,11 @@
 #include <metal_stdlib>
 
-// clang-format off
-#include "mlx/backend/metal/kernels/utils.h"
 #include "mlx/backend/metal/kernels/sdpa_vector.h"
+#include "mlx/backend/metal/kernels/utils.h"
 
 using namespace metal;
 
+// clang-format off
 // SDPA vector instantiations
 #define instantiate_sdpa_vector_aggregation(type, value_dim) \
   instantiate_kernel(                                        \
@@ -32,11 +32,9 @@ using namespace metal;
   instantiate_sdpa_vector(type, 64, 64)          \
   instantiate_sdpa_vector(type, 96, 96)          \
   instantiate_sdpa_vector(type, 128, 128)        \
-  instantiate_sdpa_vector(type, 256, 256)        \
   instantiate_sdpa_vector_aggregation(type, 64)  \
   instantiate_sdpa_vector_aggregation(type, 96)  \
-  instantiate_sdpa_vector_aggregation(type, 128) \
-  instantiate_sdpa_vector_aggregation(type, 256)
+  instantiate_sdpa_vector_aggregation(type, 128)
 
 instantiate_sdpa_vector_heads(float)
 instantiate_sdpa_vector_heads(bfloat16_t)

mlx/backend/metal/kernels/scan.h

@@ -2,8 +2,6 @@
 
 #pragma once
 
-#include "mlx/backend/metal/kernels/binary_ops.h"
-
 #define DEFINE_SIMD_SCAN()                                               \
   template <typename T, metal::enable_if_t<sizeof(T) < 8, bool> = true>  \
   T simd_scan(T val) {                                                   \
@@ -141,29 +139,6 @@ struct CumMin {
   }
 };
 
-template <typename U>
-struct CumLogaddexp {
-  static constexpr constant U init = Limits<U>::min;
-
-  template <typename T>
-  U operator()(U a, T b) {
-    return LogAddExp{}(a, static_cast<U>(b));
-  }
-
-  U simd_scan(U x) {
-    for (int i = 1; i <= 16; i *= 2) {
-      U other = simd_shuffle_and_fill_up(x, init, i);
-      x = LogAddExp{}(x, other);
-    }
-    return x;
-  }
-
-  U simd_exclusive_scan(U x) {
-    x = simd_scan(x);
-    return simd_shuffle_and_fill_up(x, init, 1);
-  }
-};
-
 template <typename T, typename U, int N_READS, bool reverse>
 inline void load_unsafe(U values[N_READS], const device T* input) {
   if (reverse) {

mlx/backend/metal/kernels/scan.metal

@@ -101,7 +101,4 @@ instantiate_scan_helper(min_int64_int64,         int64_t,     int64_t,     CumMi
 instantiate_scan_helper(min_float16_float16,     half,        half,        CumMin, 4)
 instantiate_scan_helper(min_float32_float32,     float,       float,       CumMin, 4)
 instantiate_scan_helper(min_bfloat16_bfloat16,   bfloat16_t,  bfloat16_t,  CumMin, 4)
-instantiate_scan_helper(min_complex64_complex64, complex64_t, complex64_t, CumMin, 2)
-instantiate_scan_helper(logaddexp_float16_float16,     half,        half,        CumLogaddexp, 4)
-instantiate_scan_helper(logaddexp_float32_float32,     float,       float,       CumLogaddexp, 4)
-instantiate_scan_helper(logaddexp_bfloat16_bfloat16,   bfloat16_t,  bfloat16_t,  CumLogaddexp, 4) // clang-format on
+instantiate_scan_helper(min_complex64_complex64, complex64_t, complex64_t, CumMin, 2) // clang-format on

mlx/backend/metal/kernels/sdpa_vector.h

@@ -6,9 +6,6 @@ using namespace metal;
 
 constant bool has_mask [[function_constant(20)]];
 constant bool query_transposed [[function_constant(21)]];
-constant bool do_causal [[function_constant(22)]];
-constant bool bool_mask [[function_constant(23)]];
-constant bool float_mask [[function_constant(24)]];
 
 template <typename T, int D, int V = D>
 [[kernel]] void sdpa_vector(
@@ -16,21 +13,17 @@ template <typename T, int D, int V = D>
     const device T* keys [[buffer(1)]],
     const device T* values [[buffer(2)]],
     device T* out [[buffer(3)]],
-    const constant int& gqa_factor [[buffer(4)]],
-    const constant int& N [[buffer(5)]],
-    const constant size_t& k_head_stride [[buffer(6)]],
-    const constant size_t& k_seq_stride [[buffer(7)]],
-    const constant size_t& v_head_stride [[buffer(8)]],
-    const constant size_t& v_seq_stride [[buffer(9)]],
-    const constant float& scale [[buffer(10)]],
-    const device bool* bmask [[buffer(11), function_constant(bool_mask)]],
-    const device T* fmask [[buffer(12), function_constant(float_mask)]],
-    const constant int& mask_kv_seq_stride
-    [[buffer(13), function_constant(has_mask)]],
-    const constant int& mask_q_seq_stride
-    [[buffer(14), function_constant(has_mask)]],
-    const constant int& mask_head_stride
-    [[buffer(15), function_constant(has_mask)]],
+    const constant int& gqa_factor,
+    const constant int& N,
+    const constant size_t& k_head_stride,
+    const constant size_t& k_seq_stride,
+    const constant size_t& v_head_stride,
+    const constant size_t& v_seq_stride,
+    const constant float& scale,
+    const device bool* mask [[function_constant(has_mask)]],
+    const constant int& mask_kv_seq_stride [[function_constant(has_mask)]],
+    const constant int& mask_q_seq_stride [[function_constant(has_mask)]],
+    const constant int& mask_head_stride [[function_constant(has_mask)]],
     uint3 tid [[threadgroup_position_in_grid]],
     uint3 tpg [[threadgroups_per_grid]],
     uint simd_gid [[simdgroup_index_in_threadgroup]],
@@ -64,12 +57,8 @@ template <typename T, int D, int V = D>
       simd_lid * qk_per_thread;
   values += kv_head_idx * v_head_stride + simd_gid * v_seq_stride +
       simd_lid * v_per_thread;
-  if (bool_mask) {
-    bmask += head_idx * mask_head_stride + simd_gid * mask_kv_seq_stride +
-        q_seq_idx * mask_q_seq_stride;
-  }
-  if (float_mask) {
-    fmask += head_idx * mask_head_stride + simd_gid * mask_kv_seq_stride +
+  if (has_mask) {
+    mask += head_idx * mask_head_stride + simd_gid * mask_kv_seq_stride +
         q_seq_idx * mask_q_seq_stride;
   }
 
@@ -88,13 +77,7 @@ template <typename T, int D, int V = D>
 
   // For each key
   for (int i = simd_gid; i < N; i += BN) {
-    bool use_key = true;
-    if (do_causal) {
-      use_key = i <= (N - int(tpg.y) + int(q_seq_idx));
-    } else if (bool_mask) {
-      use_key = bmask[0];
-    }
-    if (use_key) {
+    if (!has_mask || mask[0]) {
       // Read the key
       for (int j = 0; j < qk_per_thread; j++) {
         k[j] = keys[j];
@@ -106,9 +89,6 @@ template <typename T, int D, int V = D>
         score += q[j] * k[j];
       }
       score = simd_sum(score);
-      if (float_mask) {
-        score += max(Limits<U>::finite_min, static_cast<U>(fmask[0]));
-      }
 
       // Update the accumulators
       U new_max = max(max_score, score);
@@ -127,11 +107,8 @@ template <typename T, int D, int V = D>
     // Move the pointers to the next kv
     keys += inner_k_stride;
     values += inner_v_stride;
-    if (bool_mask) {
-      bmask += BN * mask_kv_seq_stride;
-    }
-    if (float_mask) {
-      fmask += BN * mask_kv_seq_stride;
+    if (has_mask) {
+      mask += BN * mask_kv_seq_stride;
     }
   }
 
@@ -172,21 +149,17 @@ template <typename T, int D, int V = D>
     device float* out [[buffer(3)]],
     device float* sums [[buffer(4)]],
     device float* maxs [[buffer(5)]],
-    const constant int& gqa_factor [[buffer(6)]],
-    const constant int& N [[buffer(7)]],
-    const constant size_t& k_head_stride [[buffer(8)]],
-    const constant size_t& k_seq_stride [[buffer(9)]],
-    const constant size_t& v_head_stride [[buffer(10)]],
-    const constant size_t& v_seq_stride [[buffer(11)]],
-    const constant float& scale [[buffer(12)]],
-    const device bool* bmask [[buffer(13), function_constant(bool_mask)]],
-    const device T* fmask [[buffer(14), function_constant(float_mask)]],
-    const constant int& mask_kv_seq_stride
-    [[buffer(15), function_constant(has_mask)]],
-    const constant int& mask_q_seq_stride
-    [[buffer(16), function_constant(has_mask)]],
-    const constant int& mask_head_stride
-    [[buffer(17), function_constant(has_mask)]],
+    const constant int& gqa_factor,
+    const constant int& N,
+    const constant size_t& k_head_stride,
+    const constant size_t& k_seq_stride,
+    const constant size_t& v_head_stride,
+    const constant size_t& v_seq_stride,
+    const constant float& scale,
+    const device bool* mask [[function_constant(has_mask)]],
+    const constant int& mask_kv_seq_stride [[function_constant(has_mask)]],
+    const constant int& mask_q_seq_stride [[function_constant(has_mask)]],
+    const constant int& mask_head_stride [[function_constant(has_mask)]],
     uint3 tid [[threadgroup_position_in_grid]],
     uint3 tpg [[threadgroups_per_grid]],
     uint simd_gid [[simdgroup_index_in_threadgroup]],
@@ -224,13 +197,8 @@ template <typename T, int D, int V = D>
   values += kv_head_idx * v_head_stride +
       (block_idx * BN + simd_gid) * v_seq_stride + simd_lid * v_per_thread;
   out += o_offset * blocks * V + block_idx * V + simd_lid * v_per_thread;
-  if (bool_mask) {
-    bmask += head_idx * mask_head_stride +
-        (block_idx * BN + simd_gid) * mask_kv_seq_stride +
-        q_seq_idx * mask_q_seq_stride;
-  }
-  if (float_mask) {
-    fmask += head_idx * mask_head_stride +
+  if (has_mask) {
+    mask += head_idx * mask_head_stride +
         (block_idx * BN + simd_gid) * mask_kv_seq_stride +
         q_seq_idx * mask_q_seq_stride;
   }
@@ -250,13 +218,7 @@ template <typename T, int D, int V = D>
 
   // For each key
   for (int i = block_idx * BN + simd_gid; i < N; i += blocks * BN) {
-    bool use_key = true;
-    if (do_causal) {
-      use_key = i <= (N - int(tpg.y) + int(q_seq_idx));
-    } else if (bool_mask) {
-      use_key = bmask[0];
-    }
-    if (use_key) {
+    if (!has_mask || mask[0]) {
       // Read the key
       for (int i = 0; i < qk_per_thread; i++) {
         k[i] = keys[i];
@@ -268,9 +230,6 @@ template <typename T, int D, int V = D>
         score += q[i] * k[i];
       }
       score = simd_sum(score);
-      if (float_mask) {
-        score += fmask[0];
-      }
 
       // Update the accumulators
       U new_max = max(max_score, score);
@@ -289,11 +248,8 @@ template <typename T, int D, int V = D>
     // Move the pointers to the next kv
     keys += blocks * inner_k_stride;
     values += blocks * inner_v_stride;
-    if (bool_mask) {
-      bmask += BN * blocks * mask_kv_seq_stride;
-    }
-    if (float_mask) {
-      fmask += BN * blocks * mask_kv_seq_stride;
+    if (has_mask) {
+      mask += BN * blocks * mask_kv_seq_stride;
     }
   }
 

mlx/backend/metal/kernels/softmax.metal

@@ -9,13 +9,47 @@ using namespace metal;
 #include "mlx/backend/metal/kernels/utils.h"
 #include "mlx/backend/metal/kernels/softmax.h"
 
-#define instantiate_softmax(name, itype)                                \
-  instantiate_kernel("block_softmax_" #name, softmax_single_row, itype) \
-  instantiate_kernel("looped_softmax_" #name, softmax_looped, itype)
+#define instantiate_softmax(name, itype)                          \
+  template [[host_name("block_softmax_" #name)]] [[kernel]] void        \
+  softmax_single_row<itype>(                                      \
+      const device itype* in,                                     \
+      device itype* out,                                          \
+      constant int& axis_size,                                    \
+      uint gid [[thread_position_in_grid]],                       \
+      uint _lid [[thread_position_in_threadgroup]],               \
+      uint simd_lane_id [[thread_index_in_simdgroup]],            \
+      uint simd_group_id [[simdgroup_index_in_threadgroup]]);     \
+  template [[host_name("looped_softmax_" #name)]] [[kernel]] void \
+  softmax_looped<itype>(                                          \
+      const device itype* in,                                     \
+      device itype* out,                                          \
+      constant int& axis_size,                                    \
+      uint gid [[threadgroup_position_in_grid]],                  \
+      uint lid [[thread_position_in_threadgroup]],                \
+      uint lsize [[threads_per_threadgroup]],                     \
+      uint simd_lane_id [[thread_index_in_simdgroup]],            \
+      uint simd_group_id [[simdgroup_index_in_threadgroup]]);
 
-#define instantiate_softmax_precise(name, itype)                                       \
-  instantiate_kernel("block_softmax_precise_" #name, softmax_single_row, itype, float) \
-  instantiate_kernel("looped_softmax_precise_" #name, softmax_looped, itype, float)
+#define instantiate_softmax_precise(name, itype)                          \
+  template [[host_name("block_softmax_precise_" #name)]] [[kernel]] void        \
+  softmax_single_row<itype, float>(                                       \
+      const device itype* in,                                             \
+      device itype* out,                                                  \
+      constant int& axis_size,                                            \
+      uint gid [[thread_position_in_grid]],                               \
+      uint _lid [[thread_position_in_threadgroup]],                       \
+      uint simd_lane_id [[thread_index_in_simdgroup]],                    \
+      uint simd_group_id [[simdgroup_index_in_threadgroup]]);             \
+  template [[host_name("looped_softmax_precise_" #name)]] [[kernel]] void \
+  softmax_looped<itype, float>(                                           \
+      const device itype* in,                                             \
+      device itype* out,                                                  \
+      constant int& axis_size,                                            \
+      uint gid [[threadgroup_position_in_grid]],                          \
+      uint lid [[thread_position_in_threadgroup]],                        \
+      uint lsize [[threads_per_threadgroup]],                             \
+      uint simd_lane_id [[thread_index_in_simdgroup]],                    \
+      uint simd_group_id [[simdgroup_index_in_threadgroup]]);
 
 instantiate_softmax(float32, float)
 instantiate_softmax(float16, half)

mlx/backend/metal/kernels/steel/attn/kernels/steel_attention.h

@@ -1,4 +1,4 @@
-// Copyright © 2024-25 Apple Inc.
+// Copyright © 2024 Apple Inc.
 
 using namespace mlx::steel;
 
@@ -9,9 +9,6 @@ using namespace mlx::steel;
 constant bool align_Q [[function_constant(200)]];
 constant bool align_K [[function_constant(201)]];
 
-constant bool has_mask [[function_constant(300)]];
-constant bool do_causal [[function_constant(301)]];
-
 template <typename T>
 struct TransformScale {
   T scale;
@@ -72,7 +69,6 @@ template <
     int BD,
     int WM,
     int WN,
-    typename MaskType = float,
     typename AccumType = float>
 [[kernel, max_total_threads_per_threadgroup(WM * WN * 32)]] void attention(
     const device T* Q [[buffer(0)]],
@@ -80,8 +76,6 @@ template <
     const device T* V [[buffer(2)]],
     device T* O [[buffer(3)]],
     const constant AttnParams* params [[buffer(4)]],
-    const constant AttnMaskParams* mask_params [[buffer(5), function_constant(has_mask)]],
-    const device MaskType* mask [[buffer(6), function_constant(has_mask)]],
     uint simd_lane_id [[thread_index_in_simdgroup]],
     uint simd_group_id [[simdgroup_index_in_threadgroup]],
     uint3 tid [[threadgroup_position_in_grid]],
@@ -108,11 +102,6 @@ template <
       tidl.y * params->O_strides[1] + // Head
       tidl.x * BQ * params->O_strides[2]; // Seqeunce
 
-  if (has_mask) {
-    mask += tidl.z * mask_params->M_strides[0] + // Batch
-        tidl.y * mask_params->M_strides[1]; // Head
-  }
-
   // Prepare threadgroup memory
   constexpr short padQ = 16 / sizeof(T);
   constexpr short padK = 16 / sizeof(T);
@@ -214,7 +203,7 @@ template <
 
   // Load Q blocks apply scale
   if (!align_Q && int(tid.x) == (params->NQ_aligned)) {
-    loader_q.load_safe(short2(BD, params->qL_rem));
+    loader_q.load_safe(short2(BD, params->qL - params->NQ_aligned * BQ));
   } else {
     loader_q.load_unsafe();
   }
@@ -229,23 +218,15 @@ template <
   // Init to -Inf
   STEEL_PRAGMA_UNROLL
   for (short i = 0; i < kRowsPT; ++i) {
-    max_score[i] = Limits<AccumType>::finite_min;
-  }
-
-  int kb_lim = params->NK;
-
-  if (do_causal) {
-    int q_max = (tid.x + 1) * BQ + params->qL_off;
-    kb_lim = (q_max + BK - 1) / BK;
-    kb_lim = min(params->NK, kb_lim);
+    max_score[i] = Limits<AccumType>::min;
   }
 
   // Loop over KV seq length
-  for (int kb = 0; kb < kb_lim; kb++) {
+  for (int kb = 0; kb < params->NK; kb++) {
     // Load K block and apply scale
     threadgroup_barrier(mem_flags::mem_threadgroup);
     if (!align_K && kb == (params->NK_aligned)) {
-      loader_k.load_safe(short2(BD, params->kL_rem));
+      loader_k.load_safe(short2(BD, params->kL - params->NK_aligned * BK));
     } else {
       loader_k.load_unsafe();
     }
@@ -269,11 +250,12 @@ template <
       tile_matmad(Stile, Qtile, Ktile, Stile);
     }
 
-    // Mask out length sequence
+    // Mask out of length sequence
     if (!align_K && kb == (params->NK_aligned)) {
       using stile_t = decltype(Stile);
       using selem_t = typename stile_t::elem_type;
-      constexpr auto neg_inf = Limits<selem_t>::finite_min;
+      constexpr auto neg_inf = -metal::numeric_limits<selem_t>::infinity();
+      const short lim = params->kL - params->NK_aligned * BK;
 
       STEEL_PRAGMA_UNROLL
       for (short i = 0; i < stile_t::kTileRows; i++) {
@@ -282,7 +264,7 @@ template <
           short col_pos = sn + (j * stile_t::kFragCols);
           STEEL_PRAGMA_UNROLL
           for (short jj = 0; jj < stile_t::MMAFrag_t::kElemCols; jj++) {
-            if ((col_pos + jj) >= params->kL_rem) {
+            if ((col_pos + jj) >= lim) {
               Stile.frag_at(i, j)[jj] = neg_inf;
             }
           }
@@ -290,78 +272,11 @@ template <
       }
     }
 
-    // Mask out if causal
-    if (do_causal && kb >= (kb_lim - ((BQ + BK - 1) / BK) - int(!align_K))) {
-      using stile_t = decltype(Stile);
-      using selem_t = typename stile_t::elem_type;
-      constexpr auto neg_inf = Limits<selem_t>::finite_min;
-
-      STEEL_PRAGMA_UNROLL
-      for (short i = 0; i < stile_t::kTileRows; i++) {
-        const int row_pos =
-            tid.x * BQ + params->qL_off + tm + sm + (i * stile_t::kFragRows);
-        STEEL_PRAGMA_UNROLL
-        for (short j = 0; j < stile_t::kTileCols; j++) {
-          const int col_pos = kb * BK + sn + (j * stile_t::kFragCols);
-          STEEL_PRAGMA_UNROLL
-          for (short jj = 0; jj < stile_t::MMAFrag_t::kElemCols; jj++) {
-            if (row_pos < (col_pos + jj)) {
-              Stile.frag_at(i, j)[jj] = neg_inf;
-            }
-          }
-        }
-      }
-    }
-
-    // Other masking as needed
-    if (has_mask) {
-      using stile_t = decltype(Stile);
-      using selem_t = typename stile_t::elem_type;
-      constexpr auto neg_inf = Limits<selem_t>::finite_min;
-
-      constexpr bool is_bool = is_same_v<MaskType, bool>;
-      using melem_t = typename metal::conditional_t<is_bool, bool, selem_t>;
-
-      using MMAFrag_mask_t = BaseMMAFrag<melem_t, kFragSize, kFragSize>;
-      using frag_t = typename MMAFrag_mask_t::frag_type;
-
-      STEEL_PRAGMA_UNROLL
-      for (short i = 0; i < stile_t::kTileRows; i++) {
-        const int row_pos = tid.x * BQ + tm + sm + (i * stile_t::kFragRows);
-        STEEL_PRAGMA_UNROLL
-        for (short j = 0; j < stile_t::kTileCols; j++) {
-          const int col_pos = kb * BK + sn + (j * stile_t::kFragCols);
-
-          frag_t mfrag;
-
-          MMAFrag_mask_t::load_safe(
-              mfrag,
-              mask,
-              int(mask_params->M_strides[2]),
-              Int<1>{},
-              params->qL,
-              params->kL,
-              row_pos,
-              col_pos);
-
-          STEEL_PRAGMA_UNROLL
-          for (short jj = 0; jj < stile_t::MMAFrag_t::kElemsPerFrag; jj++) {
-            if constexpr (is_bool) {
-              Stile.frag_at(i, j)[jj] =
-                  mfrag[jj] ? Stile.frag_at(i, j)[jj] : neg_inf;
-            } else {
-              Stile.frag_at(i, j)[jj] += 1.44269504089 * selem_t(mfrag[jj]);
-            }
-          }
-        }
-      }
-    }
-
     threadgroup_barrier(mem_flags::mem_threadgroup);
 
     // Load V blocks
     if (!align_K && kb == (params->NK_aligned)) {
-      loader_v.load_safe(short2(BD, params->kL_rem));
+      loader_v.load_safe(short2(BD, params->kL - params->NK_aligned * BK));
     } else {
       loader_v.load_unsafe();
     }
@@ -452,7 +367,8 @@ template <
   O += (tm + sm) * params->O_strides[2] + sn;
 
   if (!align_Q && int(tid.x) == (params->NQ_aligned)) {
-    auto dst_tile_dims = short2(BD - sn, params->qL_rem - (tm + sm));
+    auto dst_tile_dims =
+        short2(BD - sn, params->qL - BQ * params->NQ_aligned - (tm + sm));
 
     if (dst_tile_dims.x <= 0 || dst_tile_dims.y <= 0)
       return;

mlx/backend/metal/kernels/steel/attn/kernels/steel_attention.metal

@@ -1,4 +1,4 @@
-// Copyright © 2024-25 Apple Inc.
+// Copyright © 2024 Apple Inc.
 
 // clang-format off
 #include "mlx/backend/metal/kernels/utils.h"
@@ -6,23 +6,26 @@
 #include "mlx/backend/metal/kernels/steel/attn/attn.h"
 #include "mlx/backend/metal/kernels/steel/attn/kernels/steel_attention.h"
 
-#define instantiate_attn(tname, dtype, bq, bk, bd, wm, wn, mname, mtype) \
-  instantiate_kernel(                                                    \
-      "steel_attention_" #tname "_bq" #bq "_bk" #bk "_bd" #bd            \
-      "_wm" #wm "_wn" #wn "_mask" #mname,                                \
-  attention, dtype, bq, bk, bd, wm, wn, mtype, float)
+#define instantiate_attn(tname, dtype, bq, bk, bd, wm, wn) \
+  template [[host_name("steel_attention_" #tname "_bq" #bq "_bk" #bk "_bd" #bd "_wm" #wm "_wn" #wn)]] \
+  [[kernel]] void attention<dtype, bq, bk, bd, wm, wn, float>( \
+      const device dtype* Q [[buffer(0)]], \
+      const device dtype* K [[buffer(1)]], \
+      const device dtype* V [[buffer(2)]], \
+      device dtype* O [[buffer(3)]],\
+      const constant AttnParams* params [[buffer(4)]], \
+      uint simd_lane_id [[thread_index_in_simdgroup]], \
+      uint simd_group_id [[simdgroup_index_in_threadgroup]], \
+      uint3 tid [[threadgroup_position_in_grid]], \
+      uint3 lid [[thread_position_in_threadgroup]]);
 
-#define instantiate_attn_shapes_helper(iname, itype, mname, mtype)  \
-    instantiate_attn(iname, itype, 32, 16, 128, 4, 1, mname, mtype) \
-    instantiate_attn(iname, itype, 32, 32,  80, 4, 1, mname, mtype) \
-    instantiate_attn(iname, itype, 32, 32,  64, 4, 1, mname, mtype)
+#define instantiate_attn_shapes_helper(iname, itype) \
+    instantiate_attn(iname, itype, 32, 16, 128, 4, 1) \
+    instantiate_attn(iname, itype, 32, 32,  80, 4, 1) \
+    instantiate_attn(iname, itype, 32, 32,  64, 4, 1)
 
-#define instantiate_attn_mask_helper(iname, itype) \
-    instantiate_attn_shapes_helper(iname, itype, iname, itype) \
-    instantiate_attn_shapes_helper(iname, itype, bool_, bool)
+instantiate_attn_shapes_helper(float16, half);
+instantiate_attn_shapes_helper(bfloat16, bfloat16_t);
 
-instantiate_attn_mask_helper(float16, half);
-instantiate_attn_mask_helper(bfloat16, bfloat16_t);
-
-instantiate_attn_mask_helper(float32, float);
+instantiate_attn_shapes_helper(float32, float);
 // clang-format on

mlx/backend/metal/kernels/steel/attn/mma.h

@@ -111,7 +111,7 @@ struct BaseMMAFrag<T, 8, 8> {
       for (short j = 0; j < kElemCols; j++) {
         if ((off_x + i) < lim_x && (off_y + j) < lim_y) {
           dst[i * kElemCols + j] =
-              static_cast<T>(src[(off_x + i) * str_x + (off_y + j) * str_y]);
+              static_cast<T>(src[(off_x + i) * str_x + (off_x + j) * str_y]);
         } else {
           dst[i * kElemCols + j] = T(0);
         }

mlx/backend/metal/kernels/steel/attn/params.h

@@ -26,19 +26,11 @@ struct AttnParams {
   int NQ_aligned; ///< Number of full query blocks
   int NK_aligned; ///< Number of full key/value blocks
 
-  int qL_rem; ///< Remainder in last query block
-  int kL_rem; ///< Remainder in last key/value block
-  int qL_off; ///< Offset in query sequence start
-
   int64_t Q_strides[3]; ///< Query  strides (B, H, L, D = 1)
   int64_t K_strides[3]; ///< Key    strides (B, H, L, D = 1)
   int64_t V_strides[3]; ///< Value  strides (B, H, L, D = 1)
   int64_t O_strides[3]; ///< Output strides (B, H, L, D = 1)
 };
 
-struct AttnMaskParams {
-  int64_t M_strides[3]; ///< Mask  strides (B, H, qL, kL = 1)
-};
-
 } // namespace steel
 } // namespace mlx

mlx/backend/metal/kernels/steel/gemm/kernels/steel_gemm_fused.h

@@ -15,6 +15,10 @@ constant bool align_M [[function_constant(200)]];
 constant bool align_N [[function_constant(201)]];
 constant bool align_K [[function_constant(202)]];
 
+constant bool do_gather [[function_constant(300)]];
+
+constant bool gather_bias = do_gather && use_out_source;
+
 // clang-format off
 template <
     typename T,
@@ -35,6 +39,12 @@ template <
     const constant GEMMAddMMParams* addmm_params [[buffer(5), function_constant(use_out_source)]],
     const constant int* batch_shape [[buffer(6)]],
     const constant int64_t* batch_strides [[buffer(7)]],
+    const constant uint32_t* lhs_indices [[buffer(10), function_constant(do_gather)]],
+    const constant uint32_t* rhs_indices [[buffer(11), function_constant(do_gather)]],
+    const constant uint32_t* C_indices [[buffer(12), function_constant(gather_bias)]],
+    const constant int* operand_shape [[buffer(13), function_constant(do_gather)]],
+    const constant int64_t* operand_strides [[buffer(14), function_constant(do_gather)]],
+    const constant packed_int3& operand_batch_ndim [[buffer(15), function_constant(do_gather)]],
     uint simd_lane_id [[thread_index_in_simdgroup]],
     uint simd_group_id [[simdgroup_index_in_threadgroup]],
     uint3 tid [[threadgroup_position_in_grid]],
@@ -71,26 +81,84 @@ template <
   }
 
   // Adjust for batch
-  if (has_batch) {
-    const constant auto* A_bstrides = batch_strides;
-    const constant auto* B_bstrides = batch_strides + params->batch_ndim;
 
-    ulong2 batch_offsets = elem_to_loc_broadcast(
-        tid.z, batch_shape, A_bstrides, B_bstrides, params->batch_ndim);
+  // Handle gather
+  if (do_gather) {
+    // Read indices
+    uint32_t indx_A, indx_B, indx_C;
 
-    A += batch_offsets.x;
-    B += batch_offsets.y;
+    if (has_batch) {
+      const constant auto* indx_A_bstrides = batch_strides;
+      const constant auto* indx_B_bstrides = batch_strides + params->batch_ndim;
 
-    if (use_out_source) {
-      const constant auto* C_bstrides = B_bstrides + params->batch_ndim;
-      C += elem_to_loc(tid.z, batch_shape, C_bstrides, params->batch_ndim);
+      ulong2 indx_offsets = elem_to_loc_broadcast(
+          tid.z,
+          batch_shape,
+          indx_A_bstrides,
+          indx_B_bstrides,
+          params->batch_ndim);
+      indx_A = lhs_indices[indx_offsets.x];
+      indx_B = rhs_indices[indx_offsets.y];
+
+      if (use_out_source) {
+        const constant auto* indx_C_bstrides =
+            indx_B_bstrides + params->batch_ndim;
+        auto indx_offset_C = elem_to_loc(
+            tid.z, batch_shape, indx_C_bstrides, params->batch_ndim);
+        indx_C = C_indices[indx_offset_C];
+      }
+    } else {
+      indx_A = lhs_indices[params->batch_stride_a * tid.z];
+      indx_B = rhs_indices[params->batch_stride_b * tid.z];
+
+      if (use_out_source) {
+        indx_C = C_indices[addmm_params->batch_stride_c * tid.z];
+      }
     }
-  } else {
-    A += params->batch_stride_a * tid.z;
-    B += params->batch_stride_b * tid.z;
+
+    // Translate indices to offsets
+    int batch_ndim_A = operand_batch_ndim.x;
+    const constant int* batch_shape_A = operand_shape;
+    const constant auto* batch_strides_A = operand_strides;
+    A += elem_to_loc(indx_A, batch_shape_A, batch_strides_A, batch_ndim_A);
+
+    int batch_ndim_B = operand_batch_ndim.y;
+    const constant int* batch_shape_B = batch_shape_A + batch_ndim_A;
+    const constant auto* batch_strides_B = batch_strides_A + batch_ndim_A;
+    B += elem_to_loc(indx_B, batch_shape_B, batch_strides_B, batch_ndim_B);
 
     if (use_out_source) {
-      C += addmm_params->batch_stride_c * tid.z;
+      int batch_ndim_C = operand_batch_ndim.z;
+      const constant int* batch_shape_C = batch_shape_B + batch_ndim_B;
+      const constant auto* batch_strides_C = batch_strides_B + batch_ndim_B;
+      C += elem_to_loc(indx_C, batch_shape_C, batch_strides_C, batch_ndim_C);
+    }
+
+  }
+
+  // Handle regular batch
+  else {
+    if (has_batch) {
+      const constant auto* A_bstrides = batch_strides;
+      const constant auto* B_bstrides = batch_strides + params->batch_ndim;
+
+      ulong2 batch_offsets = elem_to_loc_broadcast(
+          tid.z, batch_shape, A_bstrides, B_bstrides, params->batch_ndim);
+
+      A += batch_offsets.x;
+      B += batch_offsets.y;
+
+      if (use_out_source) {
+        const constant auto* C_bstrides = B_bstrides + params->batch_ndim;
+        C += elem_to_loc(tid.z, batch_shape, C_bstrides, params->batch_ndim);
+      }
+    } else {
+      A += params->batch_stride_a * tid.z;
+      B += params->batch_stride_b * tid.z;
+
+      if (use_out_source) {
+        C += addmm_params->batch_stride_c * tid.z;
+      }
     }
   }
 

mlx/backend/metal/kernels/steel/gemm/kernels/steel_gemm_gather.h

@@ -1,459 +0,0 @@
-// Copyright © 2024 Apple Inc.
-
-using namespace mlx::steel;
-
-constant bool has_batch [[function_constant(10)]];
-constant bool align_M [[function_constant(200)]];
-constant bool align_N [[function_constant(201)]];
-constant bool align_K [[function_constant(202)]];
-
-template <
-    typename T,
-    int BM,
-    int BN,
-    int BK,
-    int WM,
-    int WN,
-    bool transpose_a,
-    bool transpose_b,
-    typename AccumType = float>
-[[kernel, max_total_threads_per_threadgroup(WM* WN * 32)]] void gather_mm_rhs(
-    const device T* A [[buffer(0)]],
-    const device T* B [[buffer(1)]],
-    const device uint32_t* rhs_indices [[buffer(2)]],
-    device T* C [[buffer(3)]],
-    const constant GEMMParams* params [[buffer(4)]],
-    uint simd_lane_id [[thread_index_in_simdgroup]],
-    uint simd_group_id [[simdgroup_index_in_threadgroup]],
-    uint3 tid [[threadgroup_position_in_grid]]) {
-  using gemm_kernel = GEMMKernel<
-      T,
-      T,
-      BM,
-      BN,
-      BK,
-      WM,
-      WN,
-      transpose_a,
-      transpose_b,
-      true,
-      true,
-      AccumType>;
-
-  using loader_a_t = typename gemm_kernel::loader_a_t;
-  using loader_b_t = typename gemm_kernel::loader_b_t;
-  using mma_t = typename gemm_kernel::mma_t;
-
-  if (params->tiles_n <= static_cast<int>(tid.x) ||
-      params->tiles_m <= static_cast<int>(tid.y)) {
-    return;
-  }
-
-  // Prepare threadgroup memory
-  threadgroup T As[gemm_kernel::tgp_mem_size_a];
-  threadgroup T Bs[gemm_kernel::tgp_mem_size_b];
-
-  // Find the block in A, B, C
-  const int c_row = tid.y * BM;
-  const int c_col = tid.x * BN;
-  const size_t c_row_long = size_t(c_row);
-  const size_t c_col_long = size_t(c_col);
-
-  // Prepare threadgroup bounds
-  const short tgp_bm = align_M ? BM : short(min(BM, params->M - c_row));
-  const short tgp_bn = align_N ? BN : short(min(BN, params->N - c_col));
-
-  A += transpose_a ? c_row_long : c_row_long * params->lda;
-  B += transpose_b ? c_col_long * params->ldb : c_col_long;
-  C += c_row_long * params->ldd + c_col_long;
-
-  // Do as many matmuls as necessary
-  uint32_t index;
-  short offset;
-  uint32_t index_next = rhs_indices[c_row];
-  short offset_next = 0;
-  int n = 0;
-  while (n < tgp_bm) {
-    n++;
-    offset = offset_next;
-    index = index_next;
-    offset_next = tgp_bm;
-    for (; n < tgp_bm; n++) {
-      if (rhs_indices[c_row + n] != index) {
-        offset_next = n;
-        index_next = rhs_indices[c_row + n];
-        break;
-      }
-    }
-    threadgroup_barrier(mem_flags::mem_none);
-
-    // Prepare threadgroup mma operation
-    thread mma_t mma_op(simd_group_id, simd_lane_id);
-
-    // Prepare threadgroup loading operations
-    thread loader_a_t loader_a(A, params->lda, As, simd_group_id, simd_lane_id);
-    thread loader_b_t loader_b(
-        B + index * params->batch_stride_b,
-        params->ldb,
-        Bs,
-        simd_group_id,
-        simd_lane_id);
-
-    // Prepare iterations
-    const int gemm_k_iterations = params->gemm_k_iterations_aligned;
-
-    // Do unaligned K iterations first
-    if (!align_K) {
-      const int k_last = params->gemm_k_iterations_aligned * BK;
-      const int k_remain = params->K - k_last;
-      const size_t k_jump_a =
-          transpose_a ? params->lda * size_t(k_last) : size_t(k_last);
-      const size_t k_jump_b =
-          transpose_b ? size_t(k_last) : params->ldb * size_t(k_last);
-
-      // Move loader source ahead to end
-      loader_a.src += k_jump_a;
-      loader_b.src += k_jump_b;
-
-      // Load tile
-      const short2 tile_dims_A =
-          transpose_a ? short2(tgp_bm, k_remain) : short2(k_remain, tgp_bm);
-      const short2 tile_dims_B =
-          transpose_b ? short2(k_remain, tgp_bn) : short2(tgp_bn, k_remain);
-
-      loader_a.load_safe(tile_dims_A);
-      loader_b.load_safe(tile_dims_B);
-
-      threadgroup_barrier(mem_flags::mem_threadgroup);
-
-      // Do matmul
-      mma_op.mma(As, Bs);
-
-      // Reset source back to start
-      loader_a.src -= k_jump_a;
-      loader_b.src -= k_jump_b;
-    }
-
-    // Matrix level aligned never check
-    if (align_M && align_N) {
-      for (int k = 0; k < gemm_k_iterations; k++) {
-        threadgroup_barrier(mem_flags::mem_threadgroup);
-
-        // Load elements into threadgroup
-        loader_a.load_unsafe();
-        loader_b.load_unsafe();
-
-        threadgroup_barrier(mem_flags::mem_threadgroup);
-
-        // Multiply and accumulate threadgroup elements
-        mma_op.mma(As, Bs);
-
-        // Prepare for next iteration
-        loader_a.next();
-        loader_b.next();
-      }
-
-      // Store results to device memory
-      if (offset_next - offset == BM) {
-        mma_op.store_result(C, params->ldd);
-      } else {
-        mma_op.store_result_slice(
-            C, params->ldd, short2(0, offset), short2(BN, offset_next));
-      }
-    } else {
-      const short lbk = 0;
-
-      // Tile aligned don't check
-      if ((align_M || tgp_bm == BM) && (align_N || tgp_bn == BN)) {
-        gemm_kernel::gemm_loop(
-            As,
-            Bs,
-            gemm_k_iterations,
-            loader_a,
-            loader_b,
-            mma_op,
-            tgp_bm,
-            tgp_bn,
-            lbk,
-            LoopAlignment<true, true, true>{});
-        if (offset_next - offset == BM) {
-          mma_op.store_result(C, params->ldd);
-        } else {
-          mma_op.store_result_slice(
-              C, params->ldd, short2(0, offset), short2(BN, offset_next));
-        }
-      }
-
-      // Tile partially aligned check rows
-      else if (align_N || tgp_bn == BN) {
-        gemm_kernel::gemm_loop(
-            As,
-            Bs,
-            gemm_k_iterations,
-            loader_a,
-            loader_b,
-            mma_op,
-            tgp_bm,
-            tgp_bn,
-            lbk,
-            LoopAlignment<false, true, true>{});
-        mma_op.store_result_slice(
-            C, params->ldd, short2(0, offset), short2(BN, offset_next));
-      }
-
-      // Tile partially aligned check cols
-      else if (align_M || tgp_bm == BM) {
-        gemm_kernel::gemm_loop(
-            As,
-            Bs,
-            gemm_k_iterations,
-            loader_a,
-            loader_b,
-            mma_op,
-            tgp_bm,
-            tgp_bn,
-            lbk,
-            LoopAlignment<true, false, true>{});
-        mma_op.store_result_slice(
-            C, params->ldd, short2(0, offset), short2(tgp_bn, offset_next));
-      }
-
-      // Nothing aligned so check both rows and cols
-      else {
-        gemm_kernel::gemm_loop(
-            As,
-            Bs,
-            gemm_k_iterations,
-            loader_a,
-            loader_b,
-            mma_op,
-            tgp_bm,
-            tgp_bn,
-            lbk,
-            LoopAlignment<false, false, true>{});
-        mma_op.store_result_slice(
-            C, params->ldd, short2(0, offset), short2(tgp_bn, offset_next));
-      }
-    }
-  }
-}
-
-template <
-    typename T,
-    int BM,
-    int BN,
-    int BK,
-    int WM,
-    int WN,
-    bool transpose_a,
-    bool transpose_b,
-    typename AccumType = float>
-[[kernel, max_total_threads_per_threadgroup(WM* WN * 32)]] void gather_mm(
-    const device T* A [[buffer(0)]],
-    const device T* B [[buffer(1)]],
-    const device uint32_t* lhs_indices [[buffer(2)]],
-    const device uint32_t* rhs_indices [[buffer(3)]],
-    device T* C [[buffer(4)]],
-    const constant GEMMParams* params [[buffer(5)]],
-    const constant int* indices_shape [[buffer(6)]],
-    const constant int64_t* lhs_strides [[buffer(7)]],
-    const constant int64_t* rhs_strides [[buffer(8)]],
-    const constant int& batch_ndim_a [[buffer(9)]],
-    const constant int* batch_shape_a [[buffer(10)]],
-    const constant int64_t* batch_strides_a [[buffer(11)]],
-    const constant int& batch_ndim_b [[buffer(12)]],
-    const constant int* batch_shape_b [[buffer(13)]],
-    const constant int64_t* batch_strides_b [[buffer(14)]],
-    uint simd_lane_id [[thread_index_in_simdgroup]],
-    uint simd_group_id [[simdgroup_index_in_threadgroup]],
-    uint3 tid [[threadgroup_position_in_grid]]) {
-  using gemm_kernel = GEMMKernel<
-      T,
-      T,
-      BM,
-      BN,
-      BK,
-      WM,
-      WN,
-      transpose_a,
-      transpose_b,
-      true,
-      true,
-      AccumType>;
-
-  using loader_a_t = typename gemm_kernel::loader_a_t;
-  using loader_b_t = typename gemm_kernel::loader_b_t;
-  using mma_t = typename gemm_kernel::mma_t;
-
-  if (params->tiles_n <= static_cast<int>(tid.x) ||
-      params->tiles_m <= static_cast<int>(tid.y)) {
-    return;
-  }
-
-  // Move A and B to the locations pointed by lhs_indices and rhs_indices.
-  uint32_t indx_A, indx_B;
-  if (has_batch) {
-    ulong2 indices_offsets = elem_to_loc_broadcast(
-        tid.z, indices_shape, lhs_strides, rhs_strides, params->batch_ndim);
-    indx_A = lhs_indices[indices_offsets.x];
-    indx_B = rhs_indices[indices_offsets.y];
-  } else {
-    indx_A = lhs_indices[params->batch_stride_a * tid.z];
-    indx_B = rhs_indices[params->batch_stride_b * tid.z];
-  }
-  A += elem_to_loc(indx_A, batch_shape_a, batch_strides_a, batch_ndim_a);
-  B += elem_to_loc(indx_B, batch_shape_b, batch_strides_b, batch_ndim_b);
-  C += params->batch_stride_d * tid.z;
-
-  // Prepare threadgroup memory
-  threadgroup T As[gemm_kernel::tgp_mem_size_a];
-  threadgroup T Bs[gemm_kernel::tgp_mem_size_b];
-
-  // Just make sure everybody's finished with the indexing math above.
-  threadgroup_barrier(mem_flags::mem_none);
-
-  // Find block in A, B, C
-  const int c_row = tid.y * BM;
-  const int c_col = tid.x * BN;
-  const size_t c_row_long = size_t(c_row);
-  const size_t c_col_long = size_t(c_col);
-
-  A += transpose_a ? c_row_long : c_row_long * params->lda;
-  B += transpose_b ? c_col_long * params->ldb : c_col_long;
-  C += c_row_long * params->ldd + c_col_long;
-
-  // Prepare threadgroup mma operation
-  thread mma_t mma_op(simd_group_id, simd_lane_id);
-
-  // Prepare threadgroup loading operations
-  thread loader_a_t loader_a(A, params->lda, As, simd_group_id, simd_lane_id);
-  thread loader_b_t loader_b(B, params->ldb, Bs, simd_group_id, simd_lane_id);
-
-  // Prepare threadgroup bounds
-  const short tgp_bm = align_M ? BM : short(min(BM, params->M - c_row));
-  const short tgp_bn = align_N ? BN : short(min(BN, params->N - c_col));
-
-  // Prepare iterations
-  int gemm_k_iterations = params->gemm_k_iterations_aligned;
-
-  // Do unaligned K iterations first
-  if (!align_K) {
-    const int k_last = params->gemm_k_iterations_aligned * BK;
-    const int k_remain = params->K - k_last;
-    const size_t k_jump_a =
-        transpose_a ? params->lda * size_t(k_last) : size_t(k_last);
-    const size_t k_jump_b =
-        transpose_b ? size_t(k_last) : params->ldb * size_t(k_last);
-
-    // Move loader source ahead to end
-    loader_a.src += k_jump_a;
-    loader_b.src += k_jump_b;
-
-    // Load tile
-    const short2 tile_dims_A =
-        transpose_a ? short2(tgp_bm, k_remain) : short2(k_remain, tgp_bm);
-    const short2 tile_dims_B =
-        transpose_b ? short2(k_remain, tgp_bn) : short2(tgp_bn, k_remain);
-
-    loader_a.load_safe(tile_dims_A);
-    loader_b.load_safe(tile_dims_B);
-
-    threadgroup_barrier(mem_flags::mem_threadgroup);
-
-    // Do matmul
-    mma_op.mma(As, Bs);
-
-    // Reset source back to start
-    loader_a.src -= k_jump_a;
-    loader_b.src -= k_jump_b;
-  }
-
-  // Matrix level aligned never check
-  if (align_M && align_N) {
-    for (int k = 0; k < gemm_k_iterations; k++) {
-      threadgroup_barrier(mem_flags::mem_threadgroup);
-
-      // Load elements into threadgroup
-      loader_a.load_unsafe();
-      loader_b.load_unsafe();
-
-      threadgroup_barrier(mem_flags::mem_threadgroup);
-
-      // Multiply and accumulate threadgroup elements
-      mma_op.mma(As, Bs);
-
-      // Prepare for next iteration
-      loader_a.next();
-      loader_b.next();
-    }
-
-    // Store results to device memory
-    mma_op.store_result(C, params->ldd);
-  } else {
-    const short lbk = 0;
-
-    // Tile aligned don't check
-    if ((align_M || tgp_bm == BM) && (align_N || tgp_bn == BN)) {
-      gemm_kernel::gemm_loop(
-          As,
-          Bs,
-          gemm_k_iterations,
-          loader_a,
-          loader_b,
-          mma_op,
-          tgp_bm,
-          tgp_bn,
-          lbk,
-          LoopAlignment<true, true, true>{});
-      mma_op.store_result(C, params->ldd);
-    }
-
-    // Tile partially aligned check rows
-    else if (align_N || tgp_bn == BN) {
-      gemm_kernel::gemm_loop(
-          As,
-          Bs,
-          gemm_k_iterations,
-          loader_a,
-          loader_b,
-          mma_op,
-          tgp_bm,
-          tgp_bn,
-          lbk,
-          LoopAlignment<false, true, true>{});
-      mma_op.store_result_safe(C, params->ldd, short2(tgp_bn, tgp_bm));
-    }
-
-    // Tile partially aligned check cols
-    else if (align_M || tgp_bm == BM) {
-      gemm_kernel::gemm_loop(
-          As,
-          Bs,
-          gemm_k_iterations,
-          loader_a,
-          loader_b,
-          mma_op,
-          tgp_bm,
-          tgp_bn,
-          lbk,
-          LoopAlignment<true, false, true>{});
-      mma_op.store_result_safe(C, params->ldd, short2(tgp_bn, tgp_bm));
-    }
-
-    // Nothing aligned so check both rows and cols
-    else {
-      gemm_kernel::gemm_loop(
-          As,
-          Bs,
-          gemm_k_iterations,
-          loader_a,
-          loader_b,
-          mma_op,
-          tgp_bm,
-          tgp_bn,
-          lbk,
-          LoopAlignment<false, false, true>{});
-      mma_op.store_result_safe(C, params->ldd, short2(tgp_bn, tgp_bm));
-    }
-  }
-}