MVCC: How Doximity's CUDA-to-Metal Compiler Runs Tensor Cores on M4 Macs

Hook

A compiler that performs symbolic loop analysis to recover matrix multiplication semantics from low-level CUDA instructions, then proves mathematical equivalence before generating Metal GPU code—that's not a port, that's automated theorem proving for performance.

Context

Apple silicon GPUs are powerful—the M4 Max pushes 14.2 TFLOPS for FP32—but they speak Metal, not CUDA. For ML researchers and compute engineers with existing CUDA codebases, this creates a painful choice: maintain parallel Metal implementations, run inference on CPU, or ship laptops to cloud instances with NVIDIA GPUs. PyTorch's MPS backend helps if you're already in the PyTorch ecosystem, but it won't run your custom CUDA kernels for novel attention mechanisms, quantization schemes, or sparse operations.

Doximity's mvcc (Metal Virtual CUDA Compiler) attacks this gap with a retargeting compiler that translates CUDA C++ to Metal Shading Language. Unlike traditional CUDA-to-X projects that target instruction-level compatibility, mvcc takes a radical approach: it analyzes entire loops to recover high-level matrix operations, proves the transformation preserves semantics, then generates idiomatic Metal code that maps to Apple's tensor cores. The result is a system that runs 35B MoE models at 79 tokens/second on an M4 Max—not just functional compatibility, but performance that makes Apple silicon viable for serious compute workloads.

Technical Insight

The architecture is a three-stage pipeline orchestrated by an nvcc wrapper script. First, Homebrew clang compiles your CUDA device code to LLVM IR using the NVPTX backend but without linking NVIDIA libraries—this gives you correct CUDA C++ semantics (templates, lambda captures, cooperative groups) up to the IR boundary. Second, a custom Rust translator called mvcc-ir2msl converts that LLVM IR to Metal Shading Language. Third, host code is compiled and linked against a Rust-based Metal runtime that masquerades as libcudart.dylib.

The clever part is using LLVM IR as the translation seam. By leveraging clang's existing NVPTX backend, mvcc inherits decades of compiler engineering for CUDA's C++ dialect—template instantiation, __device__ function inlining, address space inference—without reimplementing any of it. The mvcc-ir2msl translator only handles the delta between NVPTX instructions and Metal primitives. Here's what a simple kernel compilation looks like:

# Your existing CUDA code
# kernel.cu
__global__ void saxpy(int n, float a, float* x, float* y) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n) y[i] = a * x[i] + y[i];
}

# Compile with mvcc's nvcc wrapper
$ nvcc -o libkernel.dylib -shared kernel.cu
# Internally runs:
# 1. clang -target nvptx64 -emit-llvm → kernel.bc
# 2. mvcc-ir2msl kernel.bc → kernel.metal
# 3. Compile host wrapper linking libcudart.dylib

The runtime defers actual GPU code generation. When your host code calls cudaLaunchKernel for the first time, the Metal runtime hands the MSL source to Apple's compiler, compiles a pipeline state object, and caches it by hash. This costs ~1.3 seconds on first launch but means the same binary runs on M4, M5, and future Apple GPUs without recompilation—the JIT adapts to hardware capabilities at runtime.

The tensor-core recovery pass is where mvcc transcends naive translation. CUDA's mma.sync instruction exposes Tensor Cores as low-level warp-matrix operations: load 16x16 tiles with ldmatrix, accumulate with mma.m16n8k16, store with stmatrix. Translating these instruction-by-instruction to Metal SIMD operations works but yields terrible performance—you're simulating NVIDIA's ISA in software.

Instead, mvcc's recovery engine performs whole-loop symbolic analysis. It scans the IR for the characteristic pattern: nested loops with ldmatrix loads, mma.sync accumulation, and reduction epilogues. It extracts the logical GEMM dimensions (M, N, K, tile sizes), proves that the computation is equivalent to C = A @ B, and verifies that data dependencies allow vectorization. If the proof succeeds and the target GPU supports Metal 4, mvcc emits high-level cooperative tensor operations:

// Recovered from CUDA mma.sync loops
simdgroup_matrix<float, 16, 16> A, B, C;
simdgroup_load(A, src_a + row * lda, lda);
simdgroup_load(B, src_b + col, ldb);
simdgroup_matmul(C, A, B);  // Metal 4 tensor cores
simdgroup_store(C, dst + row * ldc + col, ldc);

This transformation delivers 8-21x speedups on M4/M5 because Metal 4's matmul2d operations map directly to the Apple GPU Matrix Engines. The critical safety property: if the proof fails—unusual tiling, non-affine indexing, control-flow in the loop body—mvcc falls back to exact instruction-level translation. You get correctness always, performance when provable.

Memory management reveals pragmatic compromises. cudaMalloc always allocates unified memory—there's no separate host/device memory because Apple silicon uses a unified memory architecture. This simplifies the runtime (no cudaMemcpy direction tracking, no pinned allocations) but breaks performance assumptions. CUDA code optimized for explicit transfers will see different timing profiles, and memory bandwidth is shared between CPU and GPU. The runtime implements 64-bit atomics with hashed spinlocks because Metal lacks native 64-bit atomics—a software hash table maps addresses to locks, providing correctness at the cost of potential contention on atomic-heavy workloads.

Streams map to Metal command buffers, and synchronization is eager—cudaStreamSynchronize waits for buffer completion. Warps become 32-wide SIMD groups (called simdgroups in Metal), but without independent thread scheduling. CUDA code relying on Volta-era warp divergence primitives will hit undefined behavior because Apple GPUs execute threads lockstep within a SIMD group.

Gotcha

The hard requirements cut deep: macOS 26 (unreleased as of early 2025) and M4/M5 GPUs specifically. No M1/M2/M3 support, no Intel Macs, no Linux. This is bleeding-edge silicon locked to a beta OS, which makes it a non-starter for production infrastructure or teams standardized on stable toolchains. Even if you have the hardware, you're maintaining a separate build configuration that 95% of your colleagues can't run locally.

The NVIDIA library ecosystem is completely absent. No cuBLAS, no cuDNN, no Thrust, no CUB. If your CUDA code calls cublasSgemm or uses thrust::sort, it won't link—period. This works for custom kernels written from scratch but not for PyTorch extensions, TensorFlow ops, or any code that assumes the CUDA runtime includes high-performance primitives. PyTorch's MPS backend is the better choice if you're in the Python/ML ecosystem and just need Apple GPU acceleration for standard layers.

FP64 emulation makes double-precision workloads (scientific computing, molecular dynamics, anything with double) orders of magnitude slower than native CUDA. Transcendental functions on doubles (sin, exp) are outright unsupported. Metal GPUs prioritize FP16/FP32 for graphics and ML; FP64 is an afterthought. If your workload requires double precision for numerical stability, stay on NVIDIA or move to CPU—Apple silicon will disappoint.

The tensor-core recovery is brittle by design. It only accelerates loops it can symbolically prove equivalent to matmul2d. Hand-optimized GEMM variants with swizzled memory layouts, non-standard tile sizes, or fused epilogues might fall through the analysis and hit the slow exact-translation path. The compiler won't warn you—it'll silently emit SIMD simulation code and you'll wonder why performance is 10x worse than expected. Profiling Metal frame captures becomes mandatory to verify tensor-core utilization.

Verdict

Use if: You're writing custom CUDA C++ kernels (attention mechanisms, quantization ops, sparse compute) and need them running on Apple silicon during development without maintaining parallel Metal codebases. You're on M4/M5 hardware with macOS betas, comfortable with bleeding-edge tooling, and your workloads are FP32/FP16-heavy matrix operations where the tensor-core recovery delivers big wins. The Ruby FFI integration is a sleeper hit if you're prototyping—compile CUDA at runtime, call kernels from scripts, iterate fast. Skip if: You depend on cuBLAS/cuDNN/Thrust (use PyTorch MPS instead), need stable APIs for production systems (wait for 1.0+), require FP64 scientific computing (stay on NVIDIA), or your team is on M1/M2/M3 hardware. This isn't a drop-in CUDA replacement—it's a specialized translator for teams willing to ride the cutting edge of Apple silicon for significant compilation-time savings on custom compute code.