CuMetal v0.5.0

CuMetal v0.5.0

Release index

CuMetal 0.5.0 is driven by one workload: NVIDIA Warp. Compiling it end to end
pulled in a runtime-compilation surface, the driver entry points Warp resolves
dynamically, and the CUB shim its static library needs — and, along the way,
found three miscompilations that produced wrong numbers with no diagnostic.

Highlights

• All 11 of NVIDIA Warp's libwarp CUDA sources compile through CuMetal,
  up from zero at the start of the cycle. scripts/build_warp_cumetal.sh
  reproduces it from a clean machine: it clones NVIDIA/warp at v1.12.0,
  applies the two upstream changes carried as patches in
  scripts/warp-patches/, and compiles each file, failing if any regresses.
  No fork and nothing pushed anywhere.
• Three silent miscompilations fixed. A float atomicAdd on __shared__
  memory was accepted by xcrun metallib and then dropped, so a block
  accumulator kept its initial value. Float atomics on the typed backend
  converted the payload's bit pattern numerically, so atomicAdd(p, 1.0f) added
  1065353216.0. A float selp did the same, so a comparison-selected float
  returned from a device call came back as 1082130432.0 instead of 4.0f.
  All three are the same .b32 register-typing class as the 0.2.1 bug; each now
  has a fixture that fails without its fix.
• An NVRTC and nvPTXCompiler surface, exported from libcumetal.dylib and
  aliased as libnvrtc.dylib. nvrtcCompileProgram runs cumetalc and
  nvrtcGetCUBIN returns metallib bytes that cuModuleLoadDataEx already
  accepts, so a caller written against NVRTC never learns it is driving a Metal
  toolchain. nvrtcGetPTX fails honestly rather than returning bytes it cannot
  produce.
• cuGetProcAddress, resolved against the library's own exported cu*
  symbols so it cannot drift from the real surface, plus the cuMemcpy2D,
  cuMemcpyBatchAsync, cuEventRecordWithFlags and cuStreamGetCtx entry
  points. Every driver entry point Warp resolves outside the OpenGL, IPC,
  graph-capture and CUDA-array groups is now reachable.
• cub::BlockReduce works in device code, over cub::DoubleBuffer,
  DeviceSegmentedRadixSort, and stable radix sorts. The block reduction was
  previously a host-only fallback whose temp storage could not legally live in
  __shared__.
• cudaTypedefs.h, generated from cuda.h, with the versioned PFN_cu*
  typedefs that hosts loading the driver dynamically declare; graph user
  objects; cudaGraphDebugDotPrint writing a real Graphviz file; tex1D and
  linear filtering for float2/float4 fetches.

Verification

On the recorded Apple M4 Pro system:

• Debug with the binary shim disabled passed 281/281 tests, run serially.
• Release with the binary shim disabled passed 282/282, including the
  Release-only five-kernel Phase 5 benchmark gate.
• The exact in-tree numerical corpus passed 28/28 through typed PTX and
  28/28 through direct native AOT.
• The production-metallib backend matrix passed 31/31 for direct .cu and
  typed PTX, 29/31 for the legacy PTX backend.
• The enrolled headless cuda-samples sweep, PhysX GRB conformance, llama.cpp,
  installed-tool checks, PTX sweeps, and documentation consistency all passed.

These totals are evidence for the enrolled test surfaces, not a percentage of
the complete CUDA ecosystem.

Important limits

• The Warp result is a compile result. libwarp has not been linked, no
  Warp Python test has run on a Metal device, and the CUB device-wide
  algorithms behind sort.cu and sparse.cu are host-backed — correct over
  unified memory, not device-parallel.
• bvh.cu calls cub::BlockReduce from inside a divergent branch, putting
  barriers in divergent control flow. NVIDIA hardware tolerates it; Metal may
  not. That call site is unverified.
• A __device__ function passed by value as a callback decays to a function
  pointer and the native-AOT path rejects it as an indirect call. cumetalc
  exposes no optimization level to devirtualize it.
• Apple Silicon/macOS only; no SASS execution, multi-GPU transport, NVIDIA
  management telemetry, or graphics interop.
• Typed compiler support remains bounded for deeply nested or irregular
  aggregates, indirect or multi-result device calls, and irreducible barrier
  control flow.
• Library shims are tested subsets. NVIDIA bitstream or numerical parity is not
  claimed.

See the changelog for the complete list.

CuMetal v0.4.0

CuMetal v0.4.0

Release index

CuMetal 0.4.0 is a source-first correctness and robustness release. It expands
the typed compiler and native-AOT path, hardens runtime and library boundaries,
and replaces silent compatibility success with checked behavior or explicit
unsupported returns.

Highlights

• The typed compiler supports bounded nested aggregate insert/extract/update
  paths and persistent initialized module-private globals. Native source
  executables preserve embedded metallib launch metadata and module symbols.
• Device printf now implements the statically null-format return contract,
  and graph replay preserves pointer-backed pitch and channel-aware array
  geometry.
• Runtime registration and Driver loading reject malformed or unknown framed
  fatbinary entries while retaining the documented raw PTX, ELF, LZ4, and Zstd
  compatibility forms.
• AIR inspection and validation interfaces reject malformed containers and
  inconsistent metadata with focused negative-path coverage.
• cuRAND, cuFFT, cublasLt, and cuDNN validate complete tracked allocation spans,
  checked shape/layout arithmetic, lifecycle state, and supported datatype or
  generator contracts before reading or modifying data.
• The bounded cuDNN RNN path now consumes the full input feature dimension and
  validates state and scratch geometry. The bounded attention path validates
  projection-weight offsets, configured maxima, tensor spans, and output
  overlap.
• Unified-memory advisory calls and Runtime/Driver function controls validate
  allocation ranges, streams, handles, enums, attributes, and flags even where
  Metal can provide only advisory or no-op semantics.

Verification

On the recorded Apple M4 Pro release system:

• Debug with the binary shim enabled passed 278/278 tests.
• Release with the binary shim disabled passed 276/276 tests, including the
  Release-only five-kernel Phase 5 benchmark gate.
• The fixed Phase 4 manifest passed its reviewed 185-test denominator.
• All 23/23 AIR ABI tests passed.
• PhysX, llama.cpp, the complete enrolled CUDA-samples sweep, compiler backend
  matrices, typed-PTX and native-AOT corpora, installed-tool checks, PTX sweeps,
  and documentation consistency all passed.

These totals are evidence for the enrolled test surfaces, not a percentage of
the complete CUDA ecosystem.

Important limits

• Apple Silicon/macOS only; no SASS execution, multi-GPU transport, NVIDIA
  management telemetry, or graphics interop.
• Typed compiler support remains bounded for deeply nested or irregular
  aggregates, indirect or multi-result device calls, and irreducible barrier
  control flow.
• Direct PTX texture/surface instructions, broader dynamic device printf,
  and complete CUDA graph/update semantics remain incomplete.
• Library shims are tested subsets. In particular, RNN/attention training
  breadth, general datatype/layout combinations, full algorithm selection, and
  NVIDIA bitstream or numerical parity are not claimed.
• Cross-version AIR confidence still requires genuinely distinct supported
  Xcode installations and recurring external Apple-GPU verification.

What's Changed

• Cache the inline-launch ABI sidecar parse (merges #13, with follow-ups) by @Lulzx in #14
• Cache the inline-launch ABI sidecar parse instead of re-reading it every kernel launch by @anudit in #13

New Contributors

• @anudit made their first contribution in #13

Full Changelog: v0.3.0...v0.4.0

CuMetal v0.3.0

CuMetal v0.3.0

Release notes index

CuMetal 0.3.0 is a correctness and evidence release. It closes the current
enrolled CUDA-samples gate, integrates pinned software FP64 modes, hardens major
runtime/library paths, and replaces stale monolithic documentation with linked
canonical indexes.

Highlights

• All 83/83 enrolled headless NVIDIA cuda-samples pass, with zero waivers
  and zero nonpassing manifest entries. This is a bounded snapshot, not a
  general CUDA compatibility percentage.
• Pinned VF64-metal integration provides explicit fast48, wide48, and
  software ieee64 modes while preserving binary64 storage. All three pass the
  upstream CuMetal integration gate on the recorded Apple M4 Pro system.
• CUDA graph memory nodes, dynamic cooperative partitions, source texture
  descriptors, constants/globals, device printf, memory pools, system-scope
  atomics, and additional CUDA header surfaces gained focused coverage.
• cuSPARSE SpMV gained Metal execution, stream/graph ordering, structural
  selection, and explicit reduced-precision FP64 provenance. Selected FP64
  cuBLAS level-1 work also moved from scalar CPU loops to the GPU.
• The typed compiler now has an executable 23-file four-cell production matrix:
  direct typed 9/23, typed PTX 7/23, and legacy PTX 23/23. Compile
  counts are not runtime proof.
• README, specification, status, and known-gaps material now use short index
  files with focused linked chapters and explicit evidence boundaries.

New demos

• HiGHS / cuPDLP-C GPU demo: builds the pinned solver integration, runs a
  frozen afiro CPU/GPU comparison, checks objective and residual tolerances,
  and requires honest Apple-GPU provenance. wide48 and ieee64 pass the
  recorded residual gate; fast48 reaches Optimal but misses its dual-residual
  limit.
• Mittelmann LPfeas harness: fetches a frozen feasibility corpus, records
  solver/provenance results, and generates a reproducible report without
  treating CPU fallback or approximate-stub output as GPU evidence.
• A new demo index provides one entry point for Apollo, 3D
  Gaussian Splatting, the 200k-particle SPH dam break, tiny MNIST diffusion, the
  single-sample workflow, and the new HiGHS work. The first four existed before
  v0.3.0; the index is new in this release.

Important limits

• Apple Silicon/macOS only; no SASS, multi-GPU, peer execution, or graphics
  interop.
• Typed CuMetal IR has not yet reached legacy PTX coverage and is not the PTX
  default.
• Source executables still use CUDA-compatible registration and first-launch PTX
  lowering; native source AOT descriptor integration remains open.
• Library shims and external workloads are tested subsets, not full parity.

Verification

On the recorded Apple M4 Pro release system:

• Debug with the binary shim enabled passed 247/247 non-benchmark tests;
  Release with the shim disabled passed 244/244. Both totals include the
  83/83 CUDA-samples sweep, with zero skips or failures.
• The Phase 5 benchmark gate passed, local Markdown links resolved, and both
  command-line version surfaces reported 0.3.0.
• A staged Release/shim-off installation passed cumetal doctor, compiled the
  vector-add sample through the installed cumetalc, and ran it numerically on
  the Apple GPU with exact provenance.

CuMetal v0.2.1

Patch release. If you are on 0.2.0, upgrade — it carries a silent wrong-answer bug.

Fixed

Float temporaries held in .b32 registers were typed as unsigned integers. Optimized clang PTX keeps floating-point values in .b32 %rN registers, and the PTX→MSL emitter typed neg, fma, mad, abs, min, max, not, rcp, selp, mov, and the unary math intrinsics from the register spelling rather than the instruction suffix.

A kernel computing c * (out[i] - a * p[i]) + b * q[i] emitted uint vr11 = -a;, truncating every intermediate and clamping negatives to zero:

__global__ void kmix(float *out, const float *p, const float *q,
                     float a, float b, float c, int n) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n) out[i] = c * (out[i] - a * p[i]) + b * q[i];
}
// out=2, p=q=1, a=0.5, b=0.25, c=3  ->  expected 4.75, got 6.0

Wrong answers with no diagnostic, on a shape common to most optimized float kernels. The instruction suffix is now authoritative at all of those sites, as it already was for the binary operators. This is the same flaw as the 0.1.x cvt rounding-mode bug, in the handlers that audit had flagged but not reached.

Added

A tiny diffusion-model demo (demos/diffusion) — a 312,769-parameter DDPM trained on MNIST in PyTorch, then sampled entirely by hand-written CUDA kernels through CuMetal: 1000 denoising steps, 16 images in ~13 s on an M4 Pro. run.sh --check gates a forward pass against PyTorch's own output at max |cumetal - pytorch| < 2e-3 (measured 5.2e-06). This demo is what surfaced the bug above.

Verification

221/221 tests pass serially on the Debug tree with the binary shim on (1 known environment skip, conformance_llmc_gpt2fp32cu); 216/216 on the Release tree.

Full changelog: v0.2.0...v0.2.1

CuMetal v0.2.0

CuMetal v0.2.0 adds runnable demos, closes CUDA API gaps found by the first external project built on CuMetal, and fixes seven failing tests.

Demos

• Adds the Apollo demo as the front door.
• Adds a 3D Gaussian Splatting demo running industry CUDA on Apple GPUs.
• Adds a 3D SPH dam break simulation and rendering demo.

Compatibility

• Declares cudaResourceType at namespace scope so stock CUDA sources compile without qualifying every use.
• Makes the legacy __pipeline_* helpers callable from device code.
• Replaces the blanket assert no-op with a __device__ overload, so host-side assertions in .cu files keep working.
• Exposes host malloc/free/exit via cuda_runtime.h.
• Warns when a texture object is built over linear memory while CUMETAL_USE_METAL_DEVICE_ADDRESSES is off, instead of reading zeros silently.

Correctness

• Fixes __activemask() returning zero. Any PTX special register CuMetal does not lower now refuses to lower instead of reading zero.
• Implements the %lanemask_eq/le/lt/ge/gt special registers.
• Stops one .callprototype declaration from making every entry in its module unlowerable.
• Emits the .cumetal-abi sidecar on the NVVM lowering path, fixing a launch-argument overrun.
• Fixes a double kernel launch on the source path under FP64 emulation.

Build

• Restores the PhysX GRB conformance build against PhysX 5.6.1.
• Adds layered CI groundwork and LLVM compatibility.

Suite: 221/221 passing.

Projects using CuMetal

cu_vslam_rs by @jeff-hykin brings NVIDIA cuVSLAM's CUDA GPU path to Apple Silicon, running its kernels through CuMetal without rewriting them in Metal.

Full changelog: v0.1.3...v0.2.0

CuMetal v0.1.3

CuMetal v0.1.3 simplifies the installed package and doctor output.

• Removes the installed share/cumetal/examples/vectorAdd.cu copy.
• Removes the bundled-example health check and suggested commands from cumetal doctor.
• Leaves cumetal doctor focused on installation health, ending at No issues found!.
• Retains source compilation coverage through internal test fixtures.

CuMetal v0.1.2

CuMetal v0.1.2 removes spurious PTX target-feature warnings from ordinary CUDA compilation.

• Scopes PTX version features to Clang's CUDA device compilation.
• Keeps CUDA-only features out of the Apple arm64 host compilation.
• Adds regression gates that fail if the warning returns.
• Preserves successful Apple GPU compilation and execution.

CuMetal v0.1.1

CuMetal v0.1.1 makes the installation experience more polished and immediately usable.

• Adds a Flutter-style, color-aware cumetal doctor summary.
• Respects terminal color detection, NO_COLOR, and CLICOLOR_FORCE.
• Bundles vectorAdd.cu in installed packages.
• Prints a copy-pasteable command using the real installed example path.
• Verifies the installed example by compiling and running it during packaging tests.

CuMetal v0.1.0

First experimental release of CuMetal. Includes the source-first compiler/runtime, relocatable installation, cumetal doctor, cumetal run, and the end-to-end cumetalc vectorAdd.cu -o vectorAdd workflow on Apple Silicon.