doximity/mvcc

TL;DR: mvcc is a clean-room, nvcc-compatible CUDA C++ compiler and runtime that targets Apple silicon GPUs via Metal 4 — no NVIDIA software anywhere in the loop. Unchanged CUDA sources and CMake projects build and run on macOS 26, and the standout technical idea is a compiler pass that lifts tensor-core GEMM loops (mma.sync/ldmatrix/cp.async) into whole-tile Metal 4 TensorOps, hitting 19.1 TFLOP/s f16 and ~80 tok/s decode on a 35B MoE LLM on an M5 Pro.

What it is

A compiler driver (Rust) that mimics nvcc's CLI, an IR translator (mvcc-ir2msl, C++/LLVM), and reimplementations of libcudart, libcuda and libnccl in Rust on top of Metal 4. The toolkit directory is laid out like a CUDA 12.8 install so enable_language(CUDA) and find_package(CUDAToolkit) just work. Apache-2.0, explicitly clean-room: implemented from NVIDIA's published docs, no CUDA headers, libdevice, or PTX tools shipped.

How it works

Device code is compiled to LLVM IR with stock Homebrew clang's NVPTX frontend, translated to Metal Shading Language, and embedded in the host object where clang would put a fatbinary. At first kernel launch, the runtime hands the MSL to Apple's MTLDevice.newLibraryWithSource, so codegen targets whatever GPU is actually present (pipelines cached; ~1.3 s first-launch for 50 kernels, 2–5 ms cached, or precompile .metallibs). CUDA concepts map directly: blocks→threadgroups, warps→32-wide SIMD groups, streams→command buffers, cudaMalloc→unified memory.

The clever bit: tensor-core recovery

You can't map mma.sync 1:1 onto M5's Neural Accelerators — Metal 4 TensorOps operate on whole tiles, not per-instruction. So mvcc lifts the entire GEMM loop, proves (it links Z3, so this appears to be actual formal equivalence checking) that a Metal 4 matmul2d computes the same result, and emits that. Every recovered kernel ships alongside an exact per-instruction translation built on simdgroup_matrix; the runtime falls back to it when layout verification fails or the launch shape is outside the proof. mvcc explain app.cu and MVCC_TENSOR_DIAG=1 show what recovery did per kernel. This "lift, verify, retarget" pattern is more interesting than the port itself and could generalize to other accelerator mismatches.

What's there and the evidence

Full source (compiler, runtime, headers, a CUB subset, an NCCL implementation across up to 16 logical devices on the one physical GPU), a test suite, and examples. The headline demo is examples/qwen: Qwen3.5-35B-A3B MoE inference written as plain CUDA C++, decoding at 79–81 tok/s on an M5 Pro — a real workload, not a microbenchmark. Also memory bandwidth (~243 GB/s streaming) and a Ruby gem that JIT-compiles CUDA strings and calls kernels from Ruby (the README's Ruby snippet appears garbled, likely markdown mangling).

What to expect

• No NVIDIA libraries. No cuBLAS/cuDNN/cuFFT/Thrust, so PyTorch's CUDA backend and most ML frameworks are out. This is for hand-written CUDA C++.
• First release, tested against its own tests/examples; unsupported features fail at build time with named errors (a good failure mode). Requires macOS 26, Homebrew LLVM ≥21, Rust.
• Semantics gaps to know: fp64 emulated in software (transcendentals are compile errors), no independent thread scheduling (intra-warp spinlocks can hang), device printf is silent, no NVRTC or runtime PTX loading, 32 KB fast shared memory (larger spills to a device-memory pool), 64-bit atomics via hashed spinlocks, no atomicAdd(double*).
• Numbers are honest — limitations are documented in unusual detail, and the clean-room/trademark section is careful. 19 TFLOP/s f16 is far below an RTX 4090, but competitive for a laptop-class unified-memory machine, and the 48 GB unified memory is the real draw for LLM work.

Try it

brew install llvm cmake ninja rust
git clone https://github.com/doximity/mvcc.git
cd mvcc && tools/install_toolkit.sh
export PATH="$PWD/toolkit/bin:$PATH"

nvcc -O3 -std=c++17 -o app app.cu && ./app
# or run the LLM demo in examples/qwen

For CMake projects, set CMAKE_CUDA_COMPILER and CUDAToolkit_ROOT to the toolkit and build as usual. Repo: https://github.com/doximity/mvcc