Skip to content

Releases: doximity/mvcc

mvcc 1.1.1

Choose a tag to compare

@louismullie louismullie released this 19 Oct 19:02

mvcc compiles and runs CUDA C++ on Apple silicon GPUs. It pairs an nvcc-compatible compiler, which translates device code to Metal Shading Language, with a CUDA runtime built on Metal 4.

Install

brew tap doximity/mvcc https://github.com/doximity/mvcc
brew install doximity/mvcc/cuda-mvcc
nvcc --version

Or unpack mvcc-1.1.1-macos-arm64.tar.gz and follow its INSTALL.md. The binaries are not notarized, so after a browser download run xattr -dr com.apple.quarantine toolkit once.

Requires an Apple silicon Mac running macOS 26, and clang from LLVM 23: Homebrew llvm while it is LLVM 23, or llvm@23 after that.

Changes since 1.1.0

  • mvcc-ir2msl loads Homebrew Z3 through libz3.dylib instead of the versioned soname (libz3.5.1.dylib). A newer z3 keg no longer leaves nvcc unable to start.

sha256 of mvcc-1.1.1-macos-arm64.tar.gz: 82f67182a2f5116d826081271d7b03751c9d224c1612e2052eb8a47dfab1541c

mvcc 1.1.0

Choose a tag to compare

@louismullie louismullie released this 13 Oct 18:45
d66968b

mvcc compiles and runs CUDA C++ on Apple silicon GPUs. It pairs an nvcc-compatible compiler, which translates device code to Metal Shading Language, with a CUDA runtime built on Metal 4. 1.1.0 is the first release with a prebuilt toolkit and a Homebrew formula.

Install

brew tap doximity/mvcc https://github.com/doximity/mvcc
brew install doximity/mvcc/cuda-mvcc
nvcc --version

Or unpack mvcc-1.1.0-macos-arm64.tar.gz and follow its INSTALL.md. The binaries are not notarized, so after a browser download run xattr -dr com.apple.quarantine toolkit once.

Requires an Apple silicon Mac running macOS 26, and clang from LLVM 23: Homebrew llvm while it is LLVM 23, or llvm@23 after that.

Changes since 1.0.1

  • Publish the toolkit as the Homebrew formula cuda-mvcc (Formula/cuda-mvcc.rb), which keeps the toolkit in libexec and links only nvcc and mvcc. A release asset tagged v<version> updates the formula's checksum. The release workflow accepts only a vX.Y.Z tag, and the formula text is filled from that restricted alphabet. A prefix install refuses a world-writable destination, a symlink that leaves the prefix, and a non-absolute libcudart install name, and ad-hoc signs the toolchain with the hardened runtime. The compiler drops DYLD_*, CPATH, and LIBRARY_PATH before it runs clang.
  • nvcc uses clang only from the LLVM major release that mvcc-ir2msl and its device pass were built against, now recorded in version.json: llvm@<major> first, then Homebrew llvm, then PATH. It refuses any other clang, including one named by MVCC_CLANG, with an error naming both versions. nvcc -v prints TOP and NVVMIR_LIBRARY_DIR, so CMake finds the toolkit through a symlinked nvcc.
  • cuRAND Philox4_32_10: device API in curand_kernel.h, and a host generator (curand.h, libcurand) for that generator only. Generator handles are random, and libcudart's install name is an absolute path.
  • Tensor recovery maps mma.sync.m16n8k32 with s8/u8 inputs (s32 accumulate) onto int8 TensorOps. Fragments built from plain 32-bit loads of device memory or a shared-memory tile (no ldmatrix) become matmul operands read in place, including at per-warp tile origins and through pointers in a by-value struct argument. tests/kernels/gemm_dev.cu covers bf16 and s8 GEMMs and an attention-shaped kernel.

sha256 of mvcc-1.1.0-macos-arm64.tar.gz: 0fc5e105d58341f6f4f4541a4e4ce4f41360a6fe42c950264bb40c9ec75acf4f