Request MSL 4.1 from the runtime compiler on macOS 27

[magnuslundstedt]

What

get_metal_version() has no macOS 27 branch, so every mx.fast.metal_kernel JIT compile is
capped at MSL 4.0 — even on macOS 27, whose runtime compiler supports MSL 4.1. This makes
4.1-only features (notably metal::int2b_format tensor operands for MetalPerformancePrimitives
matmul2d) unreachable through MLX on an OS that executes them fine.

This PR adds the availability branch, gated at macOS/iOS/tvOS/visionOS 27, so released OSes are
unaffected.

One spelling note: metal-cpp_26.zip (the version pinned in CMakeLists) predates the
LanguageVersion4_1 enum, so the value is written with MTLLanguageVersion's documented
(major << 16) + minor encoding and a comment. Happy to switch to MTL::LanguageVersion4_1
once Apple publishes metal-cpp_27 (404 as of today) if you'd rather wait for that.

Why it's real (measured, not inferred)

On macOS 27.0 beta (build 26A5388g, M5):

• Compiling an MPP matmul2d kernel with an int2b_format weight operand through
  MTLDevice newLibraryWithSource:options: with languageVersion = 4.1 succeeds and
  executes bit-exactly against an integer reference (ternary-valued weights, exact
  representation, max_abs_diff = 0) — on a machine with no developer tools installed;
  the OS runtime compiler alone does this.
• The same source through mx.fast.metal_kernel fails with
  no type named 'int2b_format' in namespace 'metal' and reports __METAL_VERSION__ 400,
  because MLX pins the request to 4.0.
• Negative control: on macOS 26.5.1 the runtime compiler rejects the language version itself
  (invalid value 'metal4.1' in '-std=metal4.1'), so gating at 27 is correct — there is no
  4.1 anywhere below it.

With this patch (wheel built from this branch), the identical metal_kernel source compiles at
JIT on macOS 27 and reports __METAL_VERSION__ 410; behavior on macOS 26.x is unchanged
(falls through to the LanguageVersion4_0 branch, verified). The patch builds with the
released toolchain (Xcode 26.5 SDK, no beta toolchain required) — __builtin_available
resolves the macOS 27 branch at runtime.

Minimal repro (macOS 27)

import mlx.core as mx

header = """
#include <metal_stdlib>
static_assert(__METAL_VERSION__ >= 410, "JIT did not request MSL 4.1");
using w_t = metal::int2b_format;  // 4.1-only type
"""
source = "out[thread_position_in_grid.x] = inp[thread_position_in_grid.x];"

k = mx.fast.metal_kernel(name="msl41_probe", input_names=["inp"],
                         output_names=["out"], header=header, source=source)
x = mx.arange(4, dtype=mx.float32)
print(k(inputs=[x], output_shapes=[x.shape], output_dtypes=[x.dtype],
        grid=(4, 1, 1), threadgroup=(4, 1, 1))[0])
# main: compile error (int2b_format unknown / static_assert fires, version 400)
# this branch: runs, version 410

Context

Found while measuring 2-bit weight formats for LLM inference on M5: with the runtime capped at
4.0, int2b weights (half the bytes of int4b at parity throughput through the same MPP kernels
in our measurements) can only be used via pre-compiled metallibs or direct Metal API calls,
not through MLX's JIT path.

[zcbenz]

Looks good to me, thanks!