[Metal] Compile MSL source with the highest language version the OS supports

[GY-Bai]

Problem

MetalModuleNode::GetPipelineState compiles every textual MSL module (fmt == "metal") with MTLLanguageVersion2_3, unless the device supports Metal 4 (macOS/iOS 26 SDK). MSL 2.3 has no bfloat type (added in MSL 3.1) and no device atomic<float> (added in MSL 3.0). Any kernel that uses either fails to compile at load time on macOS 13-15, even though the OS supports a newer MSL version:

Fail to compile metal source: ... error: use of undeclared identifier 'bfloat'; did you mean 'float'?

The Metal codegen prints bfloat16 as bfloat (including simdgroup_bfloat8x8 fragments). MSL supplied through tvm_callback_metal_compile, or by downstream projects that reuse the TVM Metal runtime, runs into the same limit.

Change

src/backend/metal/runtime/metal_module.mm now picks the highest MSL version below 4.0 that the running OS supports. The Metal 4 path is unchanged:

OS	MSL used
macOS 14+ / iOS 17+	3.1 (adds bfloat)
macOS 13 / iOS 16	3.0 (adds device atomic<float>)
older	2.3 (unchanged)
Metal 4 device, built with the 26 SDK	4.0 (unchanged)

Each newer version is selected inside a matching @available check, and the __MAC_OS_X_VERSION_MAX_ALLOWED / __IPHONE_OS_VERSION_MAX_ALLOWED guards follow the existing TVM_METAL_HAS_MSL_4_0 pattern. Builds against older SDKs and older deployment targets keep the current behavior.

A new test, test_metal_source_compiled_with_msl_3_1 in tests/python/codegen/test_target_codegen_metal.py, uses tvm_callback_metal_compile to compile a kernel that uses bfloat and writes out __METAL_VERSION__. It then checks that the runtime compiled it as MSL 3.1 or newer. The test is skipped below macOS 14 and when there is no Metal device.

Testing

Run on an Apple M2 with macOS 15.6.1 (Command Line Tools SDK 15.5, AppleClang 17). This is not a Metal 4 setup, so the MSL 4.0 path was not exercised.

• Built apache/tvm with USE_METAL=ON, USE_LLVM=OFF, Release, Ninja: libtvm_compiler, libtvm_runtime and libtvm_runtime_metal all built. The changed file adds no warnings. It also compiles warning-free with -mmacosx-version-min=11.0 -Wunguarded-availability.
• Because LLVM was off, kernels were built with Target("metal", host="c"), exported with export_library and loaded back with load_module.
  • New test with this change: passes. A probe kernel reports __METAL_VERSION__ == 310.
  • New test against unpatched main (only libtvm_runtime_metal rebuilt): fails with use of undeclared identifier 'bfloat'.
  • test_unaligned_vectorize still passes.
  • The float16 cases of test_metal_inf_nan and test_metal_erf could not run with the C host: the generated host C code uses half. This is a limitation of the C host only. These tests run under the LLVM host in the macOS CI job.
• Not tested on macOS 13, iOS, or a Metal 4 device.
• pre-commit run --from-ref origin/main --to-ref HEAD passes, including clang-format and ruff.

The same runtime change was proposed to the TileLang TVM fork as tile-ai#78.