Skip to content

CuMetal v0.5.0

Latest
Latest

Choose a tag to compare

View all tags
@Lulzx Lulzx released this 05 Sep 11:29
· 10 commits to main since this release
v0.5.0
507cbac

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.

Assets 2
Loading