Keyboard shortcuts

Press ← or → to navigate between chapters

Press S or / to search in the book

Press ? to show this help

Press Esc to hide this help

Rust kernels for Metal

metal-oxide compiles ordinary Rust functions into Metal compute kernels. The host uses stable Rust, typed buffers and lazy launch operations. Kernel compilation uses a separate pinned nightly and Apple’s Metal compiler.

Start with installation and the first kernel. The examples in this book are included from the repository’s executable examples; their hardware tests are explicit, separate from portable tests.

Alpha releases are published on crates.io; installation names the current version. The API can change between alpha releases.

Execution evidence

The full hardware suite has execution evidence on these devices only:

GPUmacOSXcodeNative APIs
M4 Max26.226.6classic Metal and Metal 4
Apple Paravirtual device in GitHub-hosted M1 VMs15.7 and 26.616.4 and 26.6classic Metal

Capability checks describe requirements, not proof of execution on every GPU family. The Paravirtual device lacks the Metal 4 family and Apple7 matrix support. Acquire/release atomics have compiler and capability-gate tests, but hardware validation is pending. Metal Performance Primitives (MPP) is a benchmark reference; Rust kernels expose SIMD-group matrices and do not expose a tensor API.

Installation

GPU execution requires macOS 15 or later on Apple Silicon. The host runtime supports Rust 1.89 or later; this checkout pins Rust 1.99 for development. Editing kernels also requires Xcode with its Metal Toolchain and the compiler’s pinned nightly, including rustc-dev, rust-src and llvm-tools. Install Xcode and select its developer directory before setup.

Install the alpha

In an existing Rust application, install the Cargo command and matching runtime:

cargo install cargo-metal --version '=0.1.0-alpha.1' --locked
cargo metal setup
cargo add metal-oxide@=0.1.0-alpha.1

Specify the prerelease version explicitly: plain cargo install cargo-metal skips alpha versions. Keep the runtime and Cargo command at the same version. Copy the tested source from your first kernel into src/main.rs, then run cargo run. No application build script is required.

Setup installs the pinned compiler toolchain and downloads Apple’s Metal Toolchain if missing, then checks the tools and reports the Metal device.

Run a development checkout

With Rust and Git installed, run:

git clone https://github.com/rust-dd/metal-oxide.git
cd metal-oxide
cargo run -p cargo-metal --locked -- setup
cargo run -p vec-add --locked

The vec-add example verifies several sizes against CPU results, including zero and partial blocks. To use the released source, check out v0.1.0-alpha.1 before running setup and the example.

Install the Cargo command from the same checkout to use it in other applications:

cargo install --locked --path crates/cargo-metal
cargo metal setup
cargo metal setup --check

Use a development dependency

In an existing Cargo application, choose one dependency source:

cargo add metal-oxide --path /absolute/path/to/metal-oxide/crates/metal-oxide

Or use the repository directly:

cargo add metal-oxide --git https://github.com/rust-dd/metal-oxide.git

Keep cargo-metal and the dependency at the same repository revision. Copy the tested source from your first kernel into src/main.rs and use plain cargo run. No application build script is required. For a reproducible git dependency, pin its revision and commit Cargo.lock.

Saved artifacts

Each kernel build saves metal/<file>/<item>-<math mode>/ beside the application’s sources. Commit that directory and include it when distributing a kernel crate. An unchanged saved artifact allows dependents and hosts without the kernel toolchain to build. Editing kernels or their dependencies can make it stale and requires the kernel toolchain to regenerate it.

METAL_OXIDE_ARTIFACTS=locked cargo build --locked

This mode requires current saved artifacts and fails instead of rebuilding stale ones. Saved libraries still validate their version, ABI and device requirements.

Release-profile environment settings and existing Cargo configuration files are tracked when building kernels. After adding a previously absent .cargo/config or .cargo/config.toml that changes only the release profile, touch a kernel source file or run cargo clean so Cargo expands the kernel macro again. Cargo configuration include directives are unsupported for kernel builds; put kernel settings directly in a discovered configuration file.

Kernel builds track ambient values read by env! and option_env!, including values supplied by a host build script. Cargo environment defaults retain ambient override tracking; forced values are validated through the configuration. Build scripts in device crates and their dependencies remain unsupported. Fetched packages use their published saved kernels.

Your first kernel

This is the complete vec-add executable:

use metal_oxide::device::{DisjointBuffer, ReadBuffer, thread_index};

/// Adds two buffers of equal length element by element.
#[metal_oxide::kernel(requires(a.len() == out.len(), b.len() == out.len()))]
pub fn vec_add(a: ReadBuffer<f32>, b: ReadBuffer<f32>, mut out: DisjointBuffer<f32>) {
    let index = thread_index();
    if let Some(out) = out.get_mut(&index) {
        let i = index.get();
        *out = a.as_slice()[i] + b.as_slice()[i];
    }
}

#[cfg(all(target_os = "macos", target_arch = "aarch64"))]
fn main() -> Result<(), Box<dyn std::error::Error>> {
    use metal_oxide::{Device, LaunchConfig, Operation};

    let device = Device::system_default()?;
    println!("Device: {}", device.name());
    for n in [0_u32, 1, 255, 256, 257, 1_000_003] {
        let a = (0..n).map(|i| (i % 1024) as f32 * 0.25).collect::<Vec<_>>();
        let b = (0..n).map(|i| (i % 511) as f32 * -0.5).collect::<Vec<_>>();
        let input_a = device.buffer_from_slice(&a)?;
        let input_b = device.buffer_from_slice(&b)?;
        let mut output = device.buffer_zeroed::<f32>(n as usize)?;
        let config = LaunchConfig::<256>::for_elements(n)?;
        vec_add(&device, config, &input_a, &input_b, &mut output).wait()?;
        for (i, actual) in output.read().wait().iter().enumerate() {
            if *actual != a[i] + b[i] {
                return Err(format!("vec_add mismatch at n={n}, index={i}").into());
            }
        }
        println!(
            "vec_add n={n}: verified ({} blocks x {} threads)",
            config.grid.x,
            LaunchConfig::<256>::BLOCK.x
        );
    }
    Ok(())
}

#[cfg(not(all(target_os = "macos", target_arch = "aarch64")))]
fn main() {
    eprintln!("vec-add runs on macOS with Apple Silicon");
}

#[cfg(all(test, target_os = "macos", target_arch = "aarch64"))]
mod tests;

The #[kernel] macro marks a Rust entrypoint for the compiler and produces a host launch function of the same name. Its arguments begin with the device and launch configuration, followed by the kernel’s arguments. requires relations are checked before encoding; a violation returns Error::Contract.

ReadBuffer exposes a checked shared slice. DisjointBuffer::get_mut uses a thread_index() witness to give each thread exclusive access to its own element. Extra threads get None. Here the length contracts prove the input indexing safe.

The host allocates initialized buffers, constructs a configuration, waits for the launch and reads the output. Import Operation to use .wait(). The hardware tests run only when explicitly requested; ordinary workspace tests do not run them.

For an in-place update, the tested SAXPY example uses the same ownership rules:

use metal_oxide::device::{DisjointBuffer, ReadBuffer, thread_index};

/// `y = a * x + y`: each thread reads and writes its own element of `y`.
#[metal_oxide::kernel(requires(x.len() == y.len()))]
pub fn saxpy(a: f32, x: ReadBuffer<f32>, mut y: DisjointBuffer<f32>) {
    let index = thread_index();
    if let Some(y) = y.get_mut(&index) {
        *y += a * x.as_slice()[index.get()];
    }
}

#[cfg(all(target_os = "macos", target_arch = "aarch64"))]
fn main() -> Result<(), Box<dyn std::error::Error>> {
    use metal_oxide::{Device, LaunchConfig, Operation};

    let device = Device::system_default()?;
    println!("Device: {}", device.name());
    let a = 2.5_f32;
    for n in [0_u32, 1, 255, 256, 257, 1_000_003] {
        let x = (0..n).map(|i| (i % 1024) as f32 * 0.25).collect::<Vec<_>>();
        let y = (0..n).map(|i| (i % 511) as f32 * -0.5).collect::<Vec<_>>();
        let input = device.buffer_from_slice(&x)?;
        let mut output = device.buffer_from_slice(&y)?;
        let config = LaunchConfig::<256>::for_elements(n)?;
        saxpy(&device, config, a, &input, &mut output).wait()?;
        for (i, actual) in output.read().wait().iter().enumerate() {
            if *actual != a * x[i] + y[i] {
                return Err(format!("saxpy mismatch at n={n}, index={i}").into());
            }
        }
        println!("saxpy n={n}: verified");
    }
    Ok(())
}

#[cfg(not(all(target_os = "macos", target_arch = "aarch64")))]
fn main() {
    eprintln!("saxpy runs on macOS with Apple Silicon");
}

#[cfg(all(test, target_os = "macos", target_arch = "aarch64"))]
mod tests;

Run it from the checkout with cargo run -p saxpy --locked. For larger programs, #[metal_oxide::kernels] shares types and helpers in an inline module; #[metal_oxide::kernels(path = "../kernels")] imports a separate no_std kernel crate. See compiler internals for their build boundary.

Launch geometry

LaunchConfig<X, Y, Z> stores the threads per block in its type; omitted axes are one. Its runtime grid: Dim3 counts blocks. Each block is a Metal threadgroup. LaunchConfig::<256>::for_elements(257) launches two blocks of 256 threads. The kernel must guard the 255 padding threads.

DynamicLaunchConfig carries both grid and block dimensions at runtime. It shares the const configuration’s validation: nonzero block dimensions, device and pipeline limits, total threadgroup storage and nonoverflowing global dimensions. An empty grid executes no threads but still validates its bindings and geometry. Scalar kernel parameters do not set the grid automatically.

Device coordinates

Rust functionMeaningCUDA equivalent
thread_idx()Thread within its blockthreadIdx
block_idx()Block within the gridblockIdx
block_dim()Threads per blockblockDim
grid_dim()Blocks in the gridgridDim

All four return x/y/z coordinates. The usual global x index is block_idx().x * block_dim().x + thread_idx().x. thread_index() instead returns an ownership witness for the global x index. Kernels using it, including through helpers, require one-dimensional launches; other shapes are rejected.

Required shapes

#[kernel(block = (128, 1, 1))] records an exact block requirement. Its generated host function accepts the corresponding const configuration. Host const generics alone do not specialize handwritten MSL.

These tested pipeline kernels use two required shapes and a shape-independent final stage:

/// Two stages whose kernels share an affine helper and a generic scaling body.
#[metal_oxide::kernels]
pub mod stages {
    use metal_oxide::device::{ReadBuffer, WriteBuffer, block_dim, block_idx, thread_idx};

    struct Affine {
        scale: f32,
        offset: f32,
    }

    impl Affine {
        fn apply(self, value: f32) -> f32 {
            value * self.scale + self.offset
        }
    }

    unsafe fn scale<const BLOCK: u32, const SCALE: u32>(
        input: ReadBuffer<f32>,
        output: WriteBuffer<f32>,
        n: u32,
    ) {
        let i = block_idx().x * BLOCK + thread_idx().x;
        if i < n {
            let affine = Affine {
                scale: SCALE as f32,
                offset: 3.0,
            };
            // SAFETY: entrypoint callers supply n elements and one writer per index.
            unsafe { output.store_unchecked(i, affine.apply(input.load_unchecked(i))) };
        }
    }

    /// # Safety
    /// Buffers cover n elements, do not overlap, and the launch uses 128 threads per block.
    #[kernel(block = (128, 1, 1))]
    pub unsafe fn scale_128(input: ReadBuffer<f32>, output: WriteBuffer<f32>, n: u32) {
        // SAFETY: the entrypoint contract fixes the index calculation and buffer bounds.
        unsafe { scale::<128, 2>(input, output, n) };
    }

    /// # Safety
    /// Buffers cover n elements, do not overlap, and the launch uses 256 threads per block.
    #[kernel(block = (256, 1, 1))]
    pub unsafe fn scale_256(input: ReadBuffer<f32>, output: WriteBuffer<f32>, n: u32) {
        // SAFETY: the entrypoint contract fixes the index calculation and buffer bounds.
        unsafe { scale::<256, 4>(input, output, n) };
    }

    /// # Safety
    /// Buffers cover n elements, do not overlap, and the launch is one dimensional.
    #[kernel]
    pub unsafe fn finish(input: ReadBuffer<f32>, output: WriteBuffer<f32>, n: u32) {
        let i = block_idx().x * block_dim().x + thread_idx().x;
        if i < n {
            let affine = Affine {
                scale: 5.0,
                offset: 7.0,
            };
            // SAFETY: each active thread reads an initialized input and writes its own output.
            unsafe { output.store_unchecked(i, affine.apply(input.load_unchecked(i))) };
        }
    }
}

#[cfg(all(target_os = "macos", target_arch = "aarch64"))]
mod run;

#[cfg(all(target_os = "macos", target_arch = "aarch64"))]
fn main() -> Result<(), Box<dyn std::error::Error>> {
    run::block_on(run::verify())
}

#[cfg(not(all(target_os = "macos", target_arch = "aarch64")))]
fn main() {
    eprintln!("pipeline runs on macOS with Apple Silicon");
}

#[cfg(all(test, target_os = "macos", target_arch = "aarch64"))]
mod tests;

The async chapter includes their host launches. SIMD-group matrix kernels require a declared shape whose thread count is a multiple of 32. Barriers and collectives also impose participation rules; see memory and safety.

Memory and safety

Buffer<T> owns initialized shared Metal storage. buffer_from_slice uploads values and buffer_zeroed creates initialized zero values through GpuValue. The allocation’s element count is distinct from the number of launched threads. Allocation sizes and device limits are checked.

GpuValue encodes and decodes the canonical GPU layout. Built-in primitive scalars can use direct CPU views; custom types use a typed CPU mirror. Matching Rust size or #[repr(C)] alone does not authorize interpreting arbitrary bytes as a Rust value. A #[kernels] module generates codecs for its #[repr(C)] records used in the kernel interface.

Host access and ownership

buffer.read().await returns a shared CPU view after pending GPU access completes. buffer.write().await requires an exclusive borrow and returns an exclusive view. Both also support .wait(). Check the launch result for GPU failures: CPU views wait for access but do not report command errors.

Read launch arguments accept shared borrows, ownership and Arc. Write arguments require exclusive borrows or ownership. Operations retain their arguments through execution. Devices, pipelines and buffers can cross threads; each operation’s thread traits also depend on its arguments and closures.

Within a sequence, a CPU read panics while a write is recorded. CPU writes, upload and native export panic while any access is recorded. Conflicting nested batches return InvalidBuffer. These recording reservations also protect the interval before native submission.

Device access

Prefer checked slices and DisjointBuffer for independent element updates. Unchecked buffer loads/stores and raw Launch::new require explicit unsafe contracts covering bounds, initialization, aliasing, race freedom and cooperative participation. An unsafe kernel produces an unsafe host launch function.

Thread-local, threadgroup and device memory are distinct address spaces. Threadgroup storage must be initialized before reading, and all participating threads must reach the required barriers. The compiler checks participation through helper calls and rejects varying paths that bypass barriers or collectives. It does not replace an unsafe caller’s bounds or lane-validity obligations.

ThreadgroupBuffer supplies static block-local storage. A DynamicThreadgroupBuffer<T> parameter takes an element count at launch; the runtime checks the combined static and dynamic storage against the device limit. Relaxed atomics do not order unrelated memory. Every concurrent access to an atomic element must be atomic; supported operations are listed in supported Rust.

Native integration

Device::from_native retains a native device and creates its own queues. Unsafe Buffer::from_native is an owning handoff of initialized shared, hazard-tracked storage on the same physical device. External work must finish first; native aliases must stay idle through CPU views and runtime accesses. Consuming Buffer::into_native waits for pending access and uploads CPU changes. Dropping a buffer or future alone does not authorize access through a native alias. Imported storage must match the element extent, alignment and canonical layout, and remain non-purgeable and non-aliasable while the runtime uses it.

Async operations

A generated launch returns a lazy Launch: construction submits no work. Await it with an executor, or import Operation and call .wait() to block. First-use library and pipeline preparation runs on a worker. Awaiting yields while that preparation, command capacity or GPU execution is pending.

Successful launches return arguments in the form passed. Owned and Arc read buffers can therefore travel through 'static futures; borrowed buffers remain borrowed until the operation resolves or is dropped. CPU readback is also async.

Sequences and dependencies

sequence(|batch| ...) puts pushed launches in one command buffer. Buffer accesses order dependent dispatches; independent dispatches may run concurrently. The runtime also orders conflicting operations across its queues. Device clones share that scheduling context; separate device contexts do not share resources implicitly.

The pipeline example verifies two stages against CPU results:

use std::{
    future::Future,
    pin::pin,
    sync::Arc,
    task::{Context, Poll, Wake, Waker},
};

use crate::stages;
use metal_oxide::{Device, LaunchConfig, sequence};

struct Unpark(std::thread::Thread);
impl Wake for Unpark {
    fn wake(self: Arc<Self>) {
        self.0.unpark();
    }
}

pub(super) fn block_on<F: Future>(future: F) -> F::Output {
    let waker = Waker::from(Arc::new(Unpark(std::thread::current())));
    let mut future = pin!(future);
    loop {
        match future.as_mut().poll(&mut Context::from_waker(&waker)) {
            Poll::Ready(result) => return result,
            Poll::Pending => std::thread::park(),
        }
    }
}

pub(super) async fn verify() -> Result<(), Box<dyn std::error::Error>> {
    let device = Device::system_default()?;
    for block in [128, 256] {
        for n in [0_u32, 1, 255, 256, 257, 1_000_003] {
            let values = (0..n).map(|i| (i % 1024) as f32 * 0.25).collect::<Vec<_>>();
            let input = device.buffer_from_slice(&values)?;
            let mut intermediate = device.buffer_zeroed::<f32>(n as usize)?;
            let mut output = device.buffer_zeroed::<f32>(n as usize)?;
            // SAFETY: both stages guard n, use disjoint allocations, and run in encoded order.
            sequence(|batch| unsafe {
                if block == 128 {
                    batch.push(stages::scale_128(
                        &device,
                        LaunchConfig::<128>::for_elements(n)?,
                        &input,
                        &mut intermediate,
                        n,
                    ))?;
                } else {
                    batch.push(stages::scale_256(
                        &device,
                        LaunchConfig::<256>::for_elements(n)?,
                        &input,
                        &mut intermediate,
                        n,
                    ))?;
                }
                batch.push(stages::finish(
                    &device,
                    LaunchConfig::<256>::for_elements(n)?,
                    &intermediate,
                    &mut output,
                    n,
                ))?;
                Ok(())
            })
            .await?;
            let scale = if block == 128 { 2.0 } else { 4.0 };
            for (i, &actual) in output.read().wait().iter().enumerate() {
                assert_eq!(
                    actual,
                    (values[i] * scale + 3.0) * 5.0 + 7.0,
                    "block={block}, n={n}, i={i}"
                );
            }
            println!(
                "pipeline block={block}, n={n}: verified on {}",
                device.name()
            );
        }
    }
    Ok(())
}

Its small executor is included because it is part of the tested source. An application may await the same operation using its existing async executor.

Composition and cancellation

Operation::then starts a fallible continuation after success. map transforms a successful result on the CPU; an error skips subsequent stages. shared() drives an operation once for multiple consumers, whose clones receive the same Arc of its result or error. If both futures::FutureExt and Operation are imported, use Operation::map(op, f) to disambiguate their methods.

Dropping an operation before submission releases its recording reservations and submits no work. Submitted GPU work continues, retaining its resources until completion; dropping the future does not block. Dropping a chain starts none of its later stages. Dropping every shared consumer stops driving the operation, while already submitted work completes.

A failed generated dispatch gates its dependents in the command buffer. A new batch that reads an unfinished writer’s output also carries its failure dependency. See debugging for the returned errors.

Debugging

Start with the launch’s Result. Error::Contract reports a failed host relation; InvalidLaunch reports geometry errors; unsupported artifact requirements fail before execution. An unproved bounds check, checked overflow, division by zero, literal panic/assertion or failed unwrap/expect can return Error::KernelFailed with the kernel, dispatch, Rust location and message. Dependent generated dispatches are gated; a batch depending on an unfinished failed writer returns Error::DependencyFailed.

The compiler must prove a check before a barrier or collective, or show that its failure is uniform for the participating threads. Unsupported operations fail compilation with Rust source diagnostics rather than silently changing semantics.

Validation and capture

After installing cargo-metal, run these from the checkout:

cargo metal run --validate -p vec-add --locked
cargo metal run --capture target/vec-add.gputrace -p vec-add --locked

Validation enables Metal API and shader validation. Capture records the first runtime device’s GPU work into a trace for Xcode; use a fresh path because Metal does not replace an existing trace. cargo metal test supports the same options; select a single hardware test when capturing tests.

Fresh local capture builds retain Rust line tables and sources in a separate debug artifact. Saved dependency artifacts can be captured but lack those Rust line tables. Device log! messages use Metal shader logging and appear on stderr; this requires the logging capability.

Inspect generated code

cargo metal inspect --emit ir -p vec-add --locked
cargo metal inspect --emit msl -p vec-add --locked
cargo metal inspect --emit air -p vec-add --locked

Inspection builds the selected package and prints its kernel libraries. Add --kernel NAME to select a kernel. IR shows the typed representation, MSL shows lowered control flow and checks, and AIR shows Apple’s compiled representation. cargo metal setup --check reports tool versions, installation problems and device limits without installing tools.

Performance

Measure the workload on the actual device. First-use compilation, host encoding, CPU views and GPU execution are separate costs. operation.timed() returns GPU time with the result; benchmark reports record unavailable timestamps as null. Use inspection to compare AIR when changing kernel hot paths.

Use sequence to amortize command submission for several operations. Keep intermediate data on the GPU and read back after the chain. Choose block sizes within pipeline and threadgroup-memory limits, and measure alternatives; a host const block shape does not by itself specialize shader code.

Numerical modes

The default precise mode disables implicit floating-point contraction. #[kernels(math = "fast")] explicitly permits fast functions, reassociation, reciprocal approximations and contraction. Arithmetic retains NaN and infinity support; fast functions have narrower accuracy/domain guarantees and zero signs are unspecified. The mode is part of the saved artifact. Compare equivalent modes and numerical error, rather than assuming that fast math improves throughput.

Repository benchmarks

From the checkout:

python3 scripts/bench.py --samples 20 --warmup 10 --output target/benchmarks/baseline.json

The driver checks results against CPU references and matching handwritten MSL. It covers vector addition, reduction, scan, histogram, particles, matrix multiply, sort, FFT, stencil and SIMD-group GEMM. It separates build, pipeline, host and GPU costs. --backends classic restricts a run to classic Metal; --cases LIST selects workloads. The GEMM needs an Apple7 family GPU, which the hosted Paravirtual device does not report; omit GEMM there.

Candidates interleave within a backend’s process; compare per-round ratios and repeat runs because other GPU activity and thermal state affect timings. Classic and Metal 4 run in separate processes. CPU reference work between samples can also leave the GPU idle. A command buffer containing too much work can exceed the GPU’s interactivity limit.

GEMM comparisons include MPS on classic and MPP matmul2d when available. MPP is a benchmark reference, with implementation selected by the OS; it is not a Rust tensor API. The M4 Max measurements did not meet a 1.20 Rust/MPP GPU-time threshold for both input types on both backends in every repeated run. No parity claim extends to other GPU generations. --mlx adds a separately installed reference; its host timings and result precision require care when comparing.

For the measurement protocol and recorded limitations, see the repository’s benchmark guide.

Supported Rust

The compiler imports typed rustc MIR for concrete entrypoints and their reachable helpers. Rust first checks types, lifetimes and borrowing. Kernels compile against core; arbitrary host Rust and the complete no_std ecosystem are not supported. GPU requirements appear in the generated capability table.

Accepted constructs

  • Scalars: bool, f32, F16, BF16, 8/16/32-bit integers, u64, usize and internal isize. Separate kernel crates can also use unstable primitive f16.
  • Integer powers, checked division/remainder, byte swaps, cmp/partial_cmp, derived ordering and signum; integer and float clamp, plus f32::copysign, signum and total_cmp. Invalid clamp bounds remain checked, including NaN bounds on f32.
  • Records, tuples, arrays, payload enums with default discriminants, Option, Result, exhaustive matches and ?. Enums are internal values, not ABI parameters.
  • References, slices and lifetime-bearing aggregates inside kernels and helpers. References retain their device, threadgroup or thread-local address space.
  • Concrete generic type/const instances, trait calls with a known implementation, closures and supported core iterators. Generic entrypoints declare their concrete instances(...); lifetime parameters need no instance argument.
  • Checked slice indexing, assertions and literal panics. Checks that cannot be proved become GPU failure checks where participation permits them.
  • Supported f32 math, including exercised libm and num-traits workloads. dependencies(...) declares no_std kernel dependencies; it does not make all of a dependency’s operations available.

Only owned ABI values and buffer handles cross a kernel entry boundary. Host-shared records use canonical layout codecs rather than arbitrary host memory layout. Unsupported layouts, conversions or MIR constructs produce diagnostics.

Atomics and cooperation

u32/i32 atomics support device and threadgroup load, store, exchange, weak compare/exchange, arithmetic, min/max and bitwise operations. f32 supports device load, store, exchange, weak compare/exchange and add/subtract; compare/exchange compares bits. u64 supports device min/max without returning the previous value. Orderings must be constant; SeqCst and unsupported operation/address-space combinations are rejected. Acquire/release requires MSL 4.1 and macOS 27 and still has no hardware validation evidence.

SIMD reductions, scans, shuffles, votes and barriers require their documented participation rules. simdgroup::Matrix<T> represents an 8x8 matrix distributed across 32 lanes, with f32/F16/BF16 types and supported mixed accumulation. Matrix kernels declare a block shape whose thread count is a multiple of 32. Operands and control must be uniform within each SIMD group; reference-read operands are treated as lane-varying. Loads/stores remain unsafe.

Within one thread, core::simd supports 2–4 lanes of f32 and 8/16/32-bit integers: construction, lane access, arithmetic, masks, comparisons, selection, reductions, casts, swizzles, saturation and bit counts lower to scalar operations. These vectors do not cross the ABI. Host files using this unstable API need nightly; separate kernel crates already use the pinned compiler toolchain.

Rejected constructs

Recursion, indirect calls through function pointers or trait objects, destructors, coroutines, f64, i64 and 128-bit integers are rejected. Dynamic strings and trait objects have no device representation; literal panic/expect messages are specialized during compilation. Reinterpreting pointer casts, unsupported enum layouts, wider SIMD vectors, gather/scatter and masked vector loads are rejected. Heap allocation, host I/O and operating-system facilities are outside the device subset. A varying branch or early return that bypasses required cooperative work is rejected even when the corresponding scalar code would otherwise compile.

Kernels always compile with the release profile, even when the host uses a debug build. Integer overflow checks are off by default. To enable them, add this to the workspace Cargo.toml:

[profile.release]
overflow-checks = true

Bounds, division and explicit panics remain checked regardless of this setting. Numerical transformations follow the explicit math mode.

Capability requirements

This table is generated from metal-oxide-artifact::CAPABILITIES, the same requirements used by artifact validation and runtime loading. Regenerate both the README and this chapter from the repository root:

cargo run -p metal-oxide-artifact --example capabilities --locked
cargo run -p metal-oxide-artifact --example capabilities --locked -- --check
Artifact featureMSLmacOSApple GPU familyMac2
float32_atomics3.0137no
uint64_atomics2.4128no
atomic_acquire_release4.1277no
int32_atomics1.0111yes
simd_groups2.0117yes
simdgroup_matrix2.3117no
dynamic_threadgroup_memory1.0111yes
logging3.2156yes
bfloat163.1146no

The compiler baseline is MSL 3.2 and macOS 15 even when a feature’s minimum is lower. The runtime accepts the listed Apple GPU family or, where marked yes, Mac2 support. Requirements follow Apple’s Metal language specification and feature tables. uint64_atomics covers the supported min/max operations only. Acquire/release gates are tested, but execution on the required OS has not been validated.

These are eligibility requirements. See the execution evidence for the actual tested hardware and supported Rust for accepted language constructs and operation restrictions.

Migrating from CUDA

Keep the launch’s grid and block meanings when porting a CUDA kernel:

CUDAmetal-oxide
__global__ functionRust function marked #[metal_oxide::kernel]
threadIdx, blockIdxthread_idx(), block_idx()
blockDim, gridDimblock_dim(), grid_dim()
<<<grid, block>>>generated launch with LaunchConfig or DynamicLaunchConfig
Device allocation and copyinitialized typed Buffer<T> and CPU views
__shared__ T values[N]threadgroup::shared::<T, N>() returning ThreadgroupBuffer<T, N>
extern __shared__ memoryDynamicThreadgroupBuffer<T> kernel parameter with a launch-time element count
__syncthreads()unsafe { threadgroup::barrier() }; orders threadgroup memory only
Device atomics such as atomicAddAtomicBuffer<T> and its supported operations and orderings
Warp lane IDsimdgroup::lane_id()
Kernel chain in a streamordered buffer dependencies, often in sequence

Metal threadgroups implement blocks; the grid counts blocks. Port the indexing formula and guard padding threads. For an independent one-dimensional map, prefer the thread_index() and DisjointBuffer pattern in your first kernel. For two or three dimensions, use the explicit coordinates and uphold each unsafe buffer-access contract.

There is no implicit CUDA stream handle in the public launch API. Operations are lazy and buffer dependencies order them across runtime queues. Await or wait for results before observing them on the CPU. sequence groups launches in one command buffer; a successful continuation can launch later work with then.

Port warp algorithms using the device SIMD-group API and its participation rules. Matrix operations use 32-lane groups and 8x8 matrices on supported Apple GPUs; CUDA tensor-core tile sizes and numerical behavior do not transfer automatically. MPP remains a benchmark reference. Check both the Rust subset and capabilities before translating an intrinsic.

All threads in a block must reach threadgroup::barrier() in the same loop iteration. It does not order device-buffer accesses; port synchronization that uses global memory separately.

Shared physical memory does not remove synchronization or ownership requirements. Use typed views and launch borrows instead of retaining unrestricted native aliases. Compare CPU results, numerical tolerances and performance on the target GPU.

Compiler internals

Rust #[kernel] / #[kernels]
          |
   rustc typed MIR
          |
   metal-oxide typed IR
          |             \
          MSL           ABI and generated launch functions
          |                       |
   Apple Metal compiler           |
          |                       |
      .metallib --------> stable Rust host runtime

The kernel marker does not translate Rust bodies. The pinned rustc_driver frontend imports concrete entrypoints and reachable monomorphized helpers, with device-target core metadata. Typed IR keeps layouts, source locations, reference address spaces and access modes independent of rustc internals. Structured lowering preserves checks and rejects unsupported operations explicitly.

The compiler proves checks where possible; remaining fallible operations report through a dispatch failure word. Generated kernels also have a gate that skips effects after an upstream failure. Participation analysis checks barriers and SIMD operations transitively. MSL #line directives preserve Rust diagnostics and debug locations. Apple’s compiler produces AIR and the final library; direct AIR generation remains deferred.

Host and artifact boundary

The host macros generate launch functions from the kernel signature, safety, contracts and block requirements. metal-oxide-build drives the compiler and stores the compiled library with its versioned ABI and input manifest. Source, dependencies, compiler/toolchain/SDK settings and numerical mode participate in artifact identity. Version and ABI checks are always enabled. Releases since 0.1.0-alpha0 use ABI_VERSION 4; an incompatible contract or artifact-format change increments it.

Inline modules share host/device types and generate codecs for shared records. Separate no_std kernel crates can share helpers and dependencies without putting rustc internals in the stable runtime dependency graph. The runtime validates artifact requirements, bindings, contracts and shape before encoding. It selects classic Metal or Metal 4 internally from OS/device capabilities, with separate synchronization and residency implementations behind the same public API.

Development verification

Portable workspace tests cover layouts, artifact validation and runtime contracts. Compiler tests run from crates/metal-oxide-compiler using its pinned nightly and --features rustc-private. They test source diagnostics, IR/MSL and host-executed MSL against a test-only shim. Explicit hardware suites execute the emitted shaders and compare with CPU and MSL references; compilation alone is not execution evidence. The contributing guide lists commands, and CONCEPT.md describes the implementation in depth.