`mlx` no longer builds from source with Xcode 27 (macOS 27.0 SDK)

[jrp2014]

Trying to build mlx from source fails to compile. I won't reproduce the full diagnostics here, as they are lengthy, but can be reproduced without a full rebuild by:

bash xcrun -sdk macosx metal -Wall -Wextra -fno-fast-math -Wno-c++17-extensions -Wno-c++20-extensions -c mlx/backend/metal/kernels/gated_delta_update_nax.metal -I . -o /tmp/gdn.air

The material is below is what my bot thinks, if its of any use, otherwise stop reading here.

---

☑️ I understand it is strictly prohibited to use AI to write issues.

[BUG] Metal build fails with Xcode 27 / macOS 27 SDK: gated_delta_update_nax.h passes address-space-qualified decltype to get_destination_cooperative_tensor

Summary

The gated-delta NAX kernel (mlx #4020) passes decltype(ct_a) and decltype(ct_b) to get_destination_cooperative_tensor at three call sites. Those types carry the thread address-space qualifier. The macOS 27 Metal Performance Primitives header now requires an unqualified tensor type there, so it rejects all five instantiations. The two sibling NAX kernels (steel/gemm/nax.h, steel/attn/nax.h) already wrap the operands in metal::remove_addrspace_t<…>, which is why only this file fails.

Describe the bug

mlx no longer builds from source with Xcode 27 (macOS 27.0 SDK). Compiling mlx/backend/metal/kernels/gated_delta_update_nax.metal fails with five no matching member function for call to 'get_destination_cooperative_tensor' errors.

The macOS 27 MetalPerformancePrimitives header (MPPTensorOpsMatMul2d.h) now constrains matmul2d::get_destination_cooperative_tensor with an enable_if that requires __is_tensor_type_v<LeftOperandType> || __is_cooperative_tensor_type_v<LeftOperandType> (and the same for the right operand). gated_delta_update_nax.h (added in #4020) passes bare decltype(ct_a) / decltype(ct_b), which deduce as thread metal::cooperative_tensor<...>; the thread address-space qualifier makes the trait false, so the candidate is rejected:

gated_delta_update_nax.h:129:32: error: no matching member function for call to 'get_destination_cooperative_tensor'
  auto ct_c = gemm_op.template get_destination_cooperative_tensor<
MPPTensorOpsMatMul2d.h:562:3: note: candidate template ignored: requirement 'false' was not satisfied
  [with LeftOperandType = thread metal::cooperative_tensor<float, ...>, RightOperandType = thread metal::cooperative_tensor<float, ...>, ElementType = float, CoordType = int]

The other two users of the same API, steel/gemm/nax.h and steel/attn/nax.h, already strip the qualifier with metal::remove_addrspace_t<decltype(ct_a)> and compile cleanly under the new SDK; only the three call sites in gated_delta_update_nax.h (lines 129, 175, 228 at 59d600b5e) are affected.

To Reproduce

No full build needed; from the repository root at 59d600b5e:

xcrun -sdk macosx metal -Wall -Wextra -fno-fast-math \
  -Wno-c++17-extensions -Wno-c++20-extensions \
  -c mlx/backend/metal/kernels/gated_delta_update_nax.metal -I . -o /tmp/gdn.air

gives 5 errors. pip install -e ".[dev]" fails the same way at [ 18%] Building gemv.air / gated_delta_update_nax.air.

Expected behavior

The kernel compiles, as it does with Xcode 26.

Proposed fix

Match the other NAX kernels at the three call sites:

auto ct_c = gemm_op.template get_destination_cooperative_tensor<
    metal::remove_addrspace_t<decltype(ct_a)>,
    metal::remove_addrspace_t<decltype(ct_b)>,
    CType>();

With those six operands wrapped, the same xcrun metal command produces a valid .air with no diagnostics. The change is a no-op for older SDKs, since remove_addrspace_t is already used by the sibling kernels.
Desktop:

• OS: macOS 27.0 (Darwin 27.0.0)
• Xcode 27.0 (27A266a), macOS SDK 27.0, Apple metal version 32023.921 (metalfe-32023.921.6)
• Apple M5 Max, 128 GB
• mlx main at 59d600b5e; last successful local build was 229f5b430 under Xcode 26

[jrp2014]

Applying the fix that the bot suggested above seems to work. It is necessary to enable MLX_ENABLE_TF32=0 on the test suite to get around numerical tolerance failures, but that seems to be the case at present on the M5 Max neural accelerators.