Feature - Metal GPU backend for macOS (Apple Silicon)

[rflechner]

[Backend] Metal GPU backend for macOS (Apple Silicon)

Thank you

Thank you for open-sourcing this project — the initiative is genuinely exciting and inspired me to use DotLLM.

---

Origin of the PR

It all started from a simple desire: running the samples on my MacBook. With no Metal backend available, the engine fell back to the CPU — which was expectedly much slower than what the Apple Silicon GPU could offer. Rather than living with that, I decided to implement the missing backend.

---

Summary

This PR adds a full GPU inference backend based on Apple Metal, enabling transformer model execution on Apple Silicon chips.

It is an end-to-end addition: from primitive Metal shaders all the way to a complete transformer forward pass with KV-cache and quantization, following the same architecture as the existing CUDA backend.

---

Status

Status:

• Tested on M4 Max
• Ministral-3-3B-Instruct-2512-Q4_K_M validated
• CPU parity tests added
• No multi-GPU support

---

What was added

DotLLM.Metal project

A new project in the solution, with:

• An Objective-C/C++ bridge (bridge.mm) exposing a flat C API for .NET P/Invoke.
• A public header (dotllm_metal.h) — stable API, same philosophy as the CUDA side.
• An Xcode project to develop and debug the shaders.
• The CUDA kernels were ported as faithfully as possible to Metal, to ease maintenance (especially for bug fixes) and to keep the project globally consistent.
• A build.sh script that compiles the .metal files into a pre-compiled .metallib archive.

Build prerequisites

Since Xcode 26, the Metal toolchain is no longer bundled with Xcode by default: it is a separately downloadable component. Before running build.sh, it must therefore be installed (otherwise xcrun metal fails with missing Metal Toolchain):

xcodebuild -downloadComponent MetalToolchain
cd native/metal && ./build.sh

Note that an Xcode update may remove this component and require re-running the command. This is also something to plan for in the future GitHub Action (see below).

Kernels ported from CUDA to Metal

Each Metal kernel is a port of its CUDA counterpart in the project (native/kernels/*.cu).

The split into .metal files mirrors the source .cu files 1-for-1 (same file name, same kernel names, same variable names and reduction structure), to make side-by-side review easy.

CUDA source	Metal file(s)	Kernel(s)
add.cu, add_f32.cu	add.metal, add_f32.metal	add_f16, add_f32, add_f32_f16
attention.cu, attention_f32.cu	attention.metal, attention_f32.metal	attention_f16, attention_f32
bias_add.cu, bias_add_f32.cu	bias_add.metal, bias_add_f32.metal	bias_add_f16, bias_add_f32
convert.cu	convert.metal	convert_f16_to_f32, convert_f32_to_f16
dequant.cu	dequant.metal	Q4_0, Q4_K, Q5_0, Q5_K, Q6_K, Q8_0
embedding.cu, embedding_f32out.cu	embedding.metal, embedding_f32out.metal	*_f16out and *_f32out variants
fused_add_rmsnorm.cu	fused_add_rmsnorm.metal	fused_add_rmsnorm_f16
per_head_rmsnorm.cu, per_head_rmsnorm_f32.cu	per_head_rmsnorm.metal, per_head_rmsnorm_f32.metal	per_head_rmsnorm_f16, per_head_rmsnorm_f32
quant_kv.cu	quant_kv.metal	quant_f16_to_q8_0, quant_f16_to_q4_0
quantized_gemv.cu, quantized_gemv_f32in.cu	quantized_gemv.metal, quantized_gemv_f32in.metal	Q5_0, Q4_K, Q5_K, Q6_K, Q8_0
rmsnorm.cu, rmsnorm_f32.cu, rmsnorm_f32in.cu	rmsnorm.metal, rmsnorm_f32.metal, rmsnorm_f32in.metal	rmsnorm_f16, rmsnorm_f32, rmsnorm_f32in_f16out
rope.cu, rope_f32.cu	rope.metal, rope_f32.metal	rope_f16, rope_f32
softmax.cu	softmax.metal	softmax_f16
swiglu.cu, swiglu_f32.cu	swiglu.metal, swiglu_f32.metal	swiglu_f16, swiglu_f32

CUDA → Metal mapping of the recurring constructs:

CUDA	Metal
blockIdx.x	[[threadgroup_position_in_grid]]
threadIdx.x	[[thread_position_in_threadgroup]]
blockDim.x	[[threads_per_threadgroup]]
__shfl_down_sync	simd_shuffle_down
__shared__	threadgroup
__syncthreads()	threadgroup_barrier

One intentional deviation, documented in the code:

• embedding_lookup_q6_k_f16out (embedding.metal) is a Metal-only extension, with no CUDA counterpart: the test model's embedding (Ministral-3B) is stored in Q6_K, a case the CUDA backend never had to handle.

Metal inference pipeline

• MetalWeights — loads weights from GGUF with several strategies
• MetalForwardState / IMetalForwardState — scratch buffers in unified memory (GPU and CPU variants)
• MetalTransformerModel — full transformer forward pass
• MetalModelLoader — GGUF loading adapted to the Metal backend
• MetalContext — lifecycle management of the shaders and the compiled pipeline

GPU KV-cache

A persistent KV-cache in Metal GPU memory (metal_kv_cache.mm), with dedicated tests.

Memory optimizations

• Zero-copy: GGUF tensors are registered directly in GPU memory via MTLBuffer with no intermediate copy
• Command buffer batching: reduces CPU/GPU overhead by grouping commands
• Vectorized loops in Fused Add+RMSNorm (half4) and SwiGLU (half2)

Integration

• --device metal CLI option to select the backend
• Metal support in the DotLLM.Sample.Server sample
• ModelRunContext as a unified return type for model loading
• The Windows build is not broken (the Metal files are cleanly excluded from the .csproj)
• Claude's local settings file (.claude/settings.local.json) was removed from git tracking: it is specific to each dev machine, so committing it would mix everyone's settings and add a risk of leaking local information (paths, models). It stays ignored via .gitignore.

---

Acknowledgements

This implementation was developed with assistance from Claude and ChatGPT, which helped accelerate learning and porting from CUDA to Metal.

---

Test environment

Tested only on an Apple M4 Max with 36 GB of RAM, and with a single model: Ministral-3-3B-Instruct-2512-Q4_K_M.

I tried other models (e.g. Gemma), but ran into errors about not-yet-implemented tensors. This is not specific to the Metal backend: I got exactly the same behavior with the Windows version (CUDA backend). It is therefore a pre-existing limitation of those architectures' support, not a regression introduced by this PR.

I was thus unable to validate other models or other quantization formats.

---

Remaining work / points of attention

The compiled Metal binaries (.metallib) are not versioned in the repository.
A GitHub Action running on a macOS runner and calling build.sh will probably be needed to compile them on each build.

---

What is not in this PR

• Multi-GPU support (out of scope — the CUDA backend doesn't implement it yet either)

---

Tests

Each kernel has unit tests comparing the Metal output against a scalar CPU reference, with an explicit numerical tolerance.
The Metal KV-cache is also covered by tests.

---

Notes for the reviewer

The code structure deliberately follows the same organization as the CUDA backend to ease side-by-side comparison. The Obj-C bridge exposes only a flat C API — same philosophy as dotllm_native.h.

The README.md in native/metal/ serves as a CUDA → Metal translation guide: it documents the mechanical substitutions used throughout the port (thread identity, warp/SIMD-group reductions, math intrinsics, shared memory, pointer arithmetic, grid-stride loops…). It is the reference to read before porting a new kernel, and it illustrates the mapping on the most complex cases (RoPE, attention, quantized GEMV).