CuMetal v0.7.0

CuMetal v0.7.0

Release index

CuMetal 0.7.0 makes the typed cumetal-ir backend the default on every
entry point and lands the performance roadmap: tensor cores from raw PTX,
warp-level hardware reductions and aggregated atomics, a faster single-pass
compiler driver, and runtime fast paths that exploit unified memory. Two
more external workloads run unmodified: CuPy 14.2 with CCCL 3.1 CUB, and
the pinned LAMMPS Kokkos GPU unit tests (716 of 716).

Highlights

• cumetal-ir is the universal default. cumetalc (.cu and .ptx),
  registration JIT, driver-API PTX JIT and NVRTC all lower through the typed
  backend first. JIT paths retry a typed refusal once through the legacy
  path and name the refusal in the warning; --backend=legacy /
  CUMETAL_PTX_BACKEND=legacy stays an explicit opt-in. The legacy PTX
  lowering is no longer the default anywhere.
• Tensor cores from raw PTX. wmma m16n16k16, mma.sync m16n8k16 and
  ldmatrix m8n8 x1/x2/x4 (including .trans) stage fragments through
  threadgroup memory into simdgroup_matrix operations, checked exactly
  against a host oracle on the GPU.
• Warp hardware paths. __shfl_xor_sync reduction trees and
  redux.sync become single simd_sum/simd_min/simd_max
  instructions. Uniform-address 32-bit atomic add/min/max are warp-
  aggregated: lanes combine in-warp and one leader issues the global
  atomic, with per-lane old values reconstructed by exclusive prefix.
• Typed device-side runtime. Device malloc/free, launch queues,
  device cudaMemcpyAsync queue records, and immutable function-table
  vtables with indirect calls all work through the typed backend. Mutable
  tables and unprovable indirect targets are refused, not guessed.
• Faster builds. One -O2 -S -emit-llvm host pass replaces the two
  host Clang passes (launch-stub extraction and @symbol rewrite happen
  in memory), host and device frontends run in parallel, and the metallib
  is embedded via .incbin/Mach-O section instead of a generated C byte
  array parsed by clang++ -O2.
• Codegen budgets. Dynamically indexed private arrays scalarize into
  select chains on the PTX path; threadgroup accesses coalesce to
  float4 and power-of-two 2D tiles get bank-conflict padding; an
  IR-level register-pressure estimator warns above 104 registers and
  launch bounds emit [[max_total_threads_per_threadgroup(N)]].
• Runtime overhead. Generation-validated thread-local MRU caches for
  allocation resolution and both kernel registries, a cudaMemcpy
  relocation-scan skip when no relocation can apply, .cumetal-abi
stat() bypass for native AOT kernels, a disk-backed
  MTLBinaryArchive PSO cache, and zero-copy page-aligned
  cudaHostRegister via newBufferWithBytesNoCopy.
• CMake/nvcc drop-in. CMAKE_CUDA_COMPILER=cumetalc works against
  the build-produced cumetal-toolkit tree, with -gencode,
  -Xcompiler, --expt-relaxed-constexpr and rdc detection.

Verification

At 279456c, the commit this release tags:

• Shim-ON inner loop (ctest -LE slow): 435 passed, 0 failed.
• Binary-shim OFF (-DCUMETAL_ENABLE_BINARY_SHIM=OFF): 432 passed,
  0 failed. The 3-test delta is the shim-only set.
• Unit suite: 73/73.

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

Important limits

• The legacy PTX lowering is retained as an explicit opt-in escape and
  as the JITs' named retry on typed refusal; it is not deleted.
• Driver-path modules whose code needs same-TU FP64/libdevice support
  text keep the legacy route, preserving FP64 emulation semantics there.
• Metal exposes no 64-bit atomic fetch, exchange or CAS form. Only
  discarded-result unsigned 64-bit min/max are native; signed, live-
  result, and all other 64-bit atomics keep the lock bank.
• wmma/mma staging adds static threadgroup memory; interaction with
  static shared-memory accounting has not been separately audited.
• Apple Silicon/macOS only. Library shims are tested subsets, and
  NVIDIA bitstream or numerical parity is not claimed.

See the changelog for the complete list.