Struct Graph 

pub struct Graph { /* private fields */ }

A pipeline builder. Each operation declares placement and actual buffer use.

Implementations

impl Graph

pub unsafe fn gpu_scoped( &mut self, bindings: &[GpuAccess], callback: impl FnMut(&[View<'_>]) -> Result<()> + Send + 'static, ) -> Result<Node>

Integrate owned native GPU execution into this graph’s hazards and producer dependencies. The callback runs on a worker, never in poll. Dependencies are resolved on the host before invoking dependent regions; graph completion is not enqueued as a wait on a native stream queue.

Safety

The callback must retain all private native resources it uses, access these views only as declared, and synchronize every stream touching them before returning. It must not let raw pointers/views escape this scope. Private storage must not be concurrently accessible outside the callback. On an uncertain failure return an error; captured resources and bindings are then quarantined together, preserving native lifetimes.

pub fn fill(&mut self, destination: BufferView, value: u8) -> Result<Node>

Fill a GPU-visible range with a byte pattern.

Examples found in repository

examples/graph_prepare.rs (line 13)

fn main() -> hrx::Result<()> {
    let runtime = Runtime::new()?;
    let buffer = runtime.allocate(4096, MemoryPlacement::GpuLocal)?;
    let stream = Stream::open()?;
    let raw = stream.allocate(4096)?;
    for count in [128, 512, 1024, 2048] {
        let begin = Instant::now();
        let mut graph = runtime.graph();
        for _ in 0..count {
            graph.fill(buffer.view(), 0)?;
        }
        let graph = graph.prepare()?;
        let inferred = begin.elapsed();
        let begin = Instant::now();
        let mut native = stream.graph()?;
        let mut last = None;
        for _ in 0..count {
            last = Some(native.fill(last.as_slice(), raw.binding(), 0)?);
        }
        let _native = native.finish()?;
        println!(
            "nodes={count} inferred_us={} explicit_us={}",
            inferred.as_micros(),
            begin.elapsed().as_micros()
        );
        graph.submit()?.wait()?;
    }
    Ok(())
}

examples/shared_roundtrip.rs (line 49)

fn main() -> Result<()> {
    let args: Vec<_> = std::env::args().collect();
    if args.len() > 2 {
        return Err(hrx::Error::Message(
            "usage: shared_roundtrip [bytes]".into(),
        ));
    }
    let bytes = args
        .get(1)
        .map(String::as_str)
        .unwrap_or("1048576")
        .parse::<usize>()
        .map_err(|e| hrx::Error::Message(e.to_string()))?;
    if bytes == 0 || bytes % 4096 != 0 || bytes > 64 * 1024 * 1024 {
        return Err(hrx::Error::Message(
            "bytes must be a multiple of 4096 in 4096..=67108864".into(),
        ));
    }
    let runtime = Runtime::new()?;
    let device = runtime.npu(0)?;
    let artifact = hrx::loom::Compiler::for_target(None, device.target())?
        .module(include_str!("../tests/kernels/copy.xdna.loom"))
        .compile(
            &hrx::loom::Specialization::new("copy")
                .with_config("copy.packets", (bytes / 4096).to_string()),
        )?;
    let binding = |bytes, access| BindingContract {
        bytes,
        access,
        alignment: 4,
        layout: "u32 passthrough".into(),
    };
    let contract = KernelContract {
        bindings: vec![binding(bytes, Access::Read), binding(bytes, Access::Write)],
        constants: vec![],
    };
    let kernel = unsafe { device.load_artifact(&artifact, 1, contract) }?;
    let a = runtime.allocate(bytes, MemoryPlacement::Shared(device.clone()))?;
    let c = runtime.allocate(bytes, MemoryPlacement::Shared(device.clone()))?;
    let result = runtime.allocate(bytes, MemoryPlacement::Shared(device.clone()))?;
    let mut graph = runtime.graph();
    graph.fill(a.view(), 0x5a)?;
    graph.npu(&kernel, &[a.view(), c.view()])?;
    graph.copy(result.view(), c.view())?;
    let graph = graph.prepare()?;
    for _ in 0..10 {
        graph.submit()?.wait()?;
        let mapped = result.map_read()?;
        if mapped.iter().any(|&b| b != 0x5a) {
            return Err(hrx::Error::Message("GPU/NPU roundtrip mismatch".into()));
        }
    }
    println!(
        "GPU -> NPU -> GPU: all {bytes} bytes correct, 10 executions, {} prepared regions",
        graph.regions()
    );
    Ok(())
}

examples/shared_bench.rs (line 52)

fn main() -> Result<()> {
    let args: Vec<_> = std::env::args().collect();
    if args.len() > 2 {
        return Err(hrx::Error::Message("usage: shared_bench [bytes]".into()));
    }
    let bytes = args
        .get(1)
        .map(String::as_str)
        .unwrap_or("1048576")
        .parse::<usize>()
        .map_err(|e| hrx::Error::Message(e.to_string()))?;
    if bytes == 0 || bytes % 4096 != 0 || bytes > 64 * 1024 * 1024 {
        return Err(hrx::Error::Message(
            "bytes must be a multiple of 4096 in 4096..=67108864".into(),
        ));
    }
    let runtime = Runtime::new()?;
    let device = runtime.npu(0)?;
    let artifact = hrx::loom::Compiler::for_target(None, device.target())?
        .module(include_str!("../tests/kernels/copy.xdna.loom"))
        .compile(
            &hrx::loom::Specialization::new("copy")
                .with_config("copy.packets", (bytes / 4096).to_string()),
        )?;
    let binding = |bytes, access| BindingContract {
        bytes,
        access,
        alignment: 4,
        layout: "passthrough".into(),
    };
    let kernel = unsafe {
        device.load_artifact(
            &artifact,
            1,
            KernelContract {
                bindings: vec![binding(bytes, Access::Read), binding(bytes, Access::Write)],
                constants: vec![],
            },
        )
    }?;
    let a = runtime.allocate(bytes, MemoryPlacement::Shared(device.clone()))?;
    let c = runtime.allocate(bytes, MemoryPlacement::Shared(device.clone()))?;
    let mut graph = runtime.graph();
    graph.fill(a.view(), 0x35)?;
    graph.npu(&kernel, &[a.view(), c.view()])?;
    let graph = graph.prepare()?;
    for _ in 0..5 {
        graph.submit()?.wait()?;
    }
    let mut samples = Vec::new();
    for _ in 0..100 {
        let start = Instant::now();
        graph.submit()?.wait()?;
        samples.push(start.elapsed().as_secs_f64());
    }
    if c.map_read()?.iter().any(|&b| b != 0x35) {
        return Err(hrx::Error::Message("incorrect benchmark output".into()));
    }
    samples.sort_by(f64::total_cmp);
    println!(
        "bytes={bytes} p50_us={:.2} p95_us={:.2}",
        percentile(&samples, 50) * 1e6,
        percentile(&samples, 95) * 1e6
    );
    Ok(())
}

examples/gemm_pipeline.rs (line 159)

fn main() -> Result<()> {
    let args: Vec<_> = std::env::args().collect();
    let args = if args.len() == 1 {
        vec![
            args[0].clone(),
            String::new(),
            "gemm".into(),
            "8".into(),
            "8".into(),
            "8".into(),
        ]
    } else {
        args
    };
    if args.len() != 6 {
        return Err(hrx::Error::Message(
            "usage: gemm_pipeline <gemm.loom> <symbol> M K N".into(),
        ));
    }
    let parse = |i: usize| {
        args[i]
            .parse::<usize>()
            .map_err(|e| hrx::Error::Message(e.to_string()))
    };
    let (m, k, n) = (parse(3)?, parse(4)?, parse(5)?);
    let extent = |a: usize, b: usize, width: usize| {
        a.checked_mul(b)
            .filter(|&count| count > 0 && count <= (1 << 30))
            .and_then(|count| count.checked_mul(width))
            .ok_or_else(|| hrx::Error::Message("invalid matrix extent".into()))
    };
    let (a_bytes, b_bytes, c_bytes) = (extent(m, k, 2)?, extent(k, n, 2)?, extent(m, n, 4)?);
    let runtime = Runtime::new()?;
    let device = runtime.npu(0)?;
    let (source, spec) = if args[1].is_empty() {
        (
            include_str!("../tests/kernels/gemm_bf16.xdna.loom").to_string(),
            hrx::loom::Specialization::new("gemm"),
        )
    } else {
        (
            std::fs::read_to_string(&args[1])?,
            hrx::loom::Specialization::new(&args[2])
                .with_config("gemm.m", m.to_string())
                .with_config("gemm.k", k.to_string())
                .with_config("gemm.n", n.to_string()),
        )
    };
    let artifact = hrx::loom::Compiler::for_target(None, device.target())?
        .module(&source)
        .compile(&spec)?;
    let npu = unsafe {
        device.load_artifact(
            &artifact,
            1,
            KernelContract {
                bindings: vec![
                    binding(a_bytes, Access::Read),
                    binding(b_bytes, Access::Read),
                    binding(c_bytes, Access::Write),
                ],
                constants: vec![],
            },
        )
    }?;
    let prep = gpu_kernel(
        &runtime,
        include_str!("../tests/kernels/euler.loom"),
        "krea2_euler",
        &[
            ("krea2.euler.grid_x", (m * k).div_ceil(256).to_string()),
            ("krea2.euler.grid_y", "1".into()),
        ],
        m * k,
        vec![
            binding(a_bytes, Access::ReadWrite),
            binding(a_bytes, Access::Read),
        ],
    )?;
    let epilogue = gpu_kernel(
        &runtime,
        include_str!("../tests/kernels/add_f32.loom"),
        "add_f32",
        &[("add.grid", (m * n).div_ceil(256).to_string())],
        m * n,
        vec![binding(c_bytes, Access::ReadWrite)],
    )?;
    let allocate = |bytes| runtime.allocate(bytes, MemoryPlacement::Shared(device.clone()));
    let weights = allocate(b_bytes)?;
    let ones = allocate(a_bytes)?;
    for buffer in [&weights, &ones] {
        for word in buffer.map_write()?.as_chunks_mut::<2>().0 {
            word.copy_from_slice(&0x3f80u16.to_le_bytes());
        }
    }
    let prepare = || -> Result<(ExecutableGraph, Buffer)> {
        let a = allocate(a_bytes)?;
        let c = allocate(c_bytes)?;
        let mut graph = runtime.graph();
        graph.fill(a.view(), 0)?;
        graph.gpu(&prep, &[a.view(), ones.view()])?;
        graph.npu(&npu, &[a.view(), weights.view(), c.view()])?;
        graph.gpu(&epilogue, &[c.view()])?;
        Ok((graph.prepare()?, c))
    };
    let (first, c0) = prepare()?;
    let (second, c1) = prepare()?;
    for _ in 0..5 {
        first.submit()?.wait()?;
        second.submit()?.wait()?;
    }
    let before = runtime.statistics();
    let mut latency = Vec::with_capacity(50);
    for _ in 0..50 {
        let start = Instant::now();
        first.submit()?.wait()?;
        latency.push(start.elapsed().as_secs_f64());
    }
    let start = Instant::now();
    for _ in 0..50 {
        let a = first.submit()?;
        let b = second.submit()?;
        a.wait()?;
        b.wait()?;
    }
    let requests_per_second = 100.0 / start.elapsed().as_secs_f64();
    for c in [&c0, &c1] {
        if c.map_read()?
            .as_chunks::<4>()
            .0
            .iter()
            .any(|v| f32::from_le_bytes(*v) != k as f32 + 1.0)
        {
            return Err(hrx::Error::Message("GEMM pipeline output mismatch".into()));
        }
    }
    let after = runtime.statistics();
    assert_eq!(before.allocations, after.allocations);
    assert_eq!(before.imports, after.imports);
    assert_eq!(after.copied_bytes, 0);
    latency.sort_by(f64::total_cmp);
    println!(
        "M={m} K={k} N={n} p50_ms={:.3} p95_ms={:.3} pipelined_requests_s={requests_per_second:.2}",
        percentile(&latency, 50) * 1e3,
        percentile(&latency, 95) * 1e3
    );
    println!("{}", serde_json::to_string(&after)?);
    Ok(())
}

pub fn copy( &mut self, destination: BufferView, source: BufferView, ) -> Result<Node>

Explicit GPU copy. Lengths must match and ranges must not overlap.

Examples found in repository

examples/shared_roundtrip.rs (line 51)

fn main() -> Result<()> {
    let args: Vec<_> = std::env::args().collect();
    if args.len() > 2 {
        return Err(hrx::Error::Message(
            "usage: shared_roundtrip [bytes]".into(),
        ));
    }
    let bytes = args
        .get(1)
        .map(String::as_str)
        .unwrap_or("1048576")
        .parse::<usize>()
        .map_err(|e| hrx::Error::Message(e.to_string()))?;
    if bytes == 0 || bytes % 4096 != 0 || bytes > 64 * 1024 * 1024 {
        return Err(hrx::Error::Message(
            "bytes must be a multiple of 4096 in 4096..=67108864".into(),
        ));
    }
    let runtime = Runtime::new()?;
    let device = runtime.npu(0)?;
    let artifact = hrx::loom::Compiler::for_target(None, device.target())?
        .module(include_str!("../tests/kernels/copy.xdna.loom"))
        .compile(
            &hrx::loom::Specialization::new("copy")
                .with_config("copy.packets", (bytes / 4096).to_string()),
        )?;
    let binding = |bytes, access| BindingContract {
        bytes,
        access,
        alignment: 4,
        layout: "u32 passthrough".into(),
    };
    let contract = KernelContract {
        bindings: vec![binding(bytes, Access::Read), binding(bytes, Access::Write)],
        constants: vec![],
    };
    let kernel = unsafe { device.load_artifact(&artifact, 1, contract) }?;
    let a = runtime.allocate(bytes, MemoryPlacement::Shared(device.clone()))?;
    let c = runtime.allocate(bytes, MemoryPlacement::Shared(device.clone()))?;
    let result = runtime.allocate(bytes, MemoryPlacement::Shared(device.clone()))?;
    let mut graph = runtime.graph();
    graph.fill(a.view(), 0x5a)?;
    graph.npu(&kernel, &[a.view(), c.view()])?;
    graph.copy(result.view(), c.view())?;
    let graph = graph.prepare()?;
    for _ in 0..10 {
        graph.submit()?.wait()?;
        let mapped = result.map_read()?;
        if mapped.iter().any(|&b| b != 0x5a) {
            return Err(hrx::Error::Message("GPU/NPU roundtrip mismatch".into()));
        }
    }
    println!(
        "GPU -> NPU -> GPU: all {bytes} bytes correct, 10 executions, {} prepared regions",
        graph.regions()
    );
    Ok(())
}

examples/gpu_npu_parallel.rs (line 189)

fn main() -> Result<()> {
    let args: Vec<_> = std::env::args().skip(1).collect();
    if args.len() > 4 || args.iter().any(|arg| arg == "--help") {
        println!(
            "usage: gpu_npu_parallel [gpu-elements=16777216] [npu-matrices=1024] [samples=21] [trace.json]"
        );
        return if args.len() <= 4 {
            Ok(())
        } else {
            Err(Error::Message("too many arguments".into()))
        };
    }
    let number = |index: usize, default: usize| -> Result<usize> {
        args.get(index).map_or(Ok(default), |value| {
            value
                .parse()
                .map_err(|_| Error::Message(format!("invalid integer: {value}")))
        })
    };
    let elements = number(0, 1 << 24)?;
    let matrices = number(1, 1024)?;
    let samples = number(2, 21)?;
    if !(256..=1 << 26).contains(&elements)
        || !elements.is_multiple_of(256)
        || !(1..=16384).contains(&matrices)
        || !(3..=1000).contains(&samples)
    {
        return Err(Error::Message("gpu-elements must be a multiple of 256 in 256..=67108864; npu-matrices in 1..=16384; samples in 3..=1000".into()));
    }
    let trace_path = args.get(3).map(PathBuf::from);
    let runtime = Runtime::new()?;
    let gpu = runtime.gpu()?;
    let npu = runtime.npu(0)?;
    println!(
        "GPU: {}; NPU: {}",
        gpu.target().as_str(),
        npu.target().as_str()
    );
    let gpu_compiler = Compiler::for_target(None, gpu.target())?;
    let gpu_artifact = gpu_compiler
        .import_cxx(CxxSource::new(
            "affine.cpp",
            format!(
                r#"
#include <hip/hip_runtime.h>
__global__ [[loom::workgroup_size(256, 1, 1), loom::workgroup_count({}, 1, 1)]]
void affine(const int* input, int* output) {{
  unsigned i = blockIdx.x * 256u + threadIdx.x;
  output[i] = input[i] * 3 + 7;
}}
"#,
                elements / 256
            ),
        ))?
        .compile(&Specialization::new("affine"))?;
    // SAFETY: the fixed launch covers exactly elements i32 values in each
    // nonaliasing binding; bounded input values cannot overflow the arithmetic.
    let gpu_kernel = unsafe {
        runtime.load_gpu_artifact(
            &gpu_artifact,
            &[],
            contract(&[(elements * 4, Access::Read), (elements * 4, Access::Write)]),
        )
    }?;
    let npu_artifact = Compiler::for_target(None, npu.target())?
        .module(include_str!("kernels/batched_gemm.xdna.loom"))
        .compile(&Specialization::new("gemm").with_config("gemm.batch", matrices.to_string()))?;
    // SAFETY: the pipeline consumes matrices pairs of 8x8 BF16 inputs and
    // produces the same number of 8x8 FP32 outputs on one NPU column.
    let npu_kernel = unsafe {
        npu.load_artifact(
            &npu_artifact,
            1,
            contract(&[
                (matrices * 128, Access::Read),
                (matrices * 128, Access::Read),
                (matrices * 256, Access::Write),
            ]),
        )
    }?;
    let input = runtime.allocate(elements * 4, MemoryPlacement::GpuLocal)?;
    let output = runtime.allocate(elements * 4, MemoryPlacement::GpuLocal)?;
    let host_input = runtime.allocate(elements * 4, MemoryPlacement::HostVisible)?;
    let host_output = runtime.allocate(elements * 4, MemoryPlacement::HostVisible)?;
    let mut upload = runtime.graph();
    upload.copy(input.view(), host_input.view())?;
    let upload = upload.prepare()?;
    let mut download = runtime.graph();
    download.copy(host_output.view(), output.view())?;
    let download = download.prepare()?;
    let a = runtime.allocate(matrices * 128, MemoryPlacement::NpuLocal(npu.clone()))?;
    let b = runtime.allocate(matrices * 128, MemoryPlacement::NpuLocal(npu.clone()))?;
    let c = runtime.allocate(matrices * 256, MemoryPlacement::NpuLocal(npu))?;

    let mut gpu_only = runtime.graph();
    gpu_only.gpu(&gpu_kernel, &[input.view(), output.view()])?;
    let gpu_only = gpu_only.prepare()?;
    let mut npu_only = runtime.graph();
    npu_only.npu(&npu_kernel, &[a.view(), b.view(), c.view()])?;
    let npu_only = npu_only.prepare()?;

    // Disjoint buffers imply no dependency edge. One submit admits both
    // branches to separate GPU/NPU workers; one wait joins their completions.
    let mut parallel = runtime.graph();
    parallel.gpu(&gpu_kernel, &[input.view(), output.view()])?;
    parallel.npu(&npu_kernel, &[a.view(), b.view(), c.view()])?;
    let parallel = parallel.prepare()?;
    for pass in 0..3 {
        initialize(&host_input, &a, &b, pass * 2)?;
        upload.submit()?.wait()?;
        gpu_only.submit()?.wait()?;
        npu_only.submit()?.wait()?;
        download.submit()?.wait()?;
        validate(&host_output, &c, pass * 2)?;
        // Change both inputs so validation cannot accept the sequential run's
        // old output if either parallel branch fails to write its result.
        initialize(&host_input, &a, &b, pass * 2 + 1)?;
        upload.submit()?.wait()?;
        parallel.submit()?.wait()?;
        download.submit()?.wait()?;
        validate(&host_output, &c, pass * 2 + 1)?;
    }

    let before = runtime.statistics();
    let mut sequential_ms = Vec::with_capacity(samples);
    let mut parallel_ms = Vec::with_capacity(samples);
    for sample in 0..samples {
        // Alternate order to reduce systematic clock/thermal ordering bias.
        for concurrent in if sample % 2 == 0 {
            [false, true]
        } else {
            [true, false]
        } {
            if concurrent {
                parallel_ms.push(completed_ms(&parallel)?);
            } else {
                let start = Instant::now();
                gpu_only.submit()?.wait()?;
                npu_only.submit()?.wait()?;
                sequential_ms.push(start.elapsed().as_secs_f64() * 1000.0);
            }
        }
    }
    let after = runtime.statistics();
    if before.allocations != after.allocations || before.imports != after.imports {
        return Err(Error::Message(
            "replay unexpectedly allocated or imported storage".into(),
        ));
    }
    download.submit()?.wait()?;
    validate(&host_output, &c, 5)?;
    let sequential = median(&mut sequential_ms);
    let concurrent = median(&mut parallel_ms);
    println!("{elements} GPU elements + {matrices} NPU matrices; {samples} paired samples");
    println!(
        "sequential median: {sequential:.3} ms; parallel median: {concurrent:.3} ms; speedup: {:.2}x",
        sequential / concurrent
    );
    println!(
        "Both CPU references match exactly, including changed-input replays. No native allocations/imports during timed replay."
    );

    // Capture separately so tracing does not affect the benchmark samples.
    runtime.start_trace(32)?;
    for _ in 0..4 {
        parallel.submit()?.wait()?;
    }
    let trace = runtime
        .finish_trace()
        .ok_or_else(|| Error::Message("missing execution trace".into()))?;
    let overlap: f64 = trace
        .events
        .iter()
        .filter(|e| e.lane == 1)
        .flat_map(|gpu| {
            trace.events.iter().filter(|e| e.lane == 3).map(move |npu| {
                ((gpu.start_us + gpu.duration_us).min(npu.start_us + npu.duration_us)
                    - gpu.start_us.max(npu.start_us))
                .max(0.0)
            })
        })
        .sum();
    println!(
        "GPU/NPU host-observed overlap across four replays: {overlap:.1} us (not device timestamps)"
    );
    if let Some(path) = trace_path {
        std::fs::write(&path, trace.to_json()?)?;
        println!("Trace: {}", path.display());
    }
    Ok(())
}

pub fn copy_on( &mut self, destination: BufferView, source: BufferView, lane: GpuLane, ) -> Result<Node>

Copy on an explicit lane, retaining the same cross-lane hazard checks.

pub fn gpu( &mut self, kernel: &GpuKernel, bindings: &[BufferView], ) -> Result<Node>

Invoke a trusted GPU kernel; binding extents and alignment are checked. The selected runtime must expose allocation-address queries (interop ABI 1), including for GPU-local bindings. No base-pointer alignment is assumed.

Examples found in repository

examples/gemm_pipeline.rs (line 160)

fn main() -> Result<()> {
    let args: Vec<_> = std::env::args().collect();
    let args = if args.len() == 1 {
        vec![
            args[0].clone(),
            String::new(),
            "gemm".into(),
            "8".into(),
            "8".into(),
            "8".into(),
        ]
    } else {
        args
    };
    if args.len() != 6 {
        return Err(hrx::Error::Message(
            "usage: gemm_pipeline <gemm.loom> <symbol> M K N".into(),
        ));
    }
    let parse = |i: usize| {
        args[i]
            .parse::<usize>()
            .map_err(|e| hrx::Error::Message(e.to_string()))
    };
    let (m, k, n) = (parse(3)?, parse(4)?, parse(5)?);
    let extent = |a: usize, b: usize, width: usize| {
        a.checked_mul(b)
            .filter(|&count| count > 0 && count <= (1 << 30))
            .and_then(|count| count.checked_mul(width))
            .ok_or_else(|| hrx::Error::Message("invalid matrix extent".into()))
    };
    let (a_bytes, b_bytes, c_bytes) = (extent(m, k, 2)?, extent(k, n, 2)?, extent(m, n, 4)?);
    let runtime = Runtime::new()?;
    let device = runtime.npu(0)?;
    let (source, spec) = if args[1].is_empty() {
        (
            include_str!("../tests/kernels/gemm_bf16.xdna.loom").to_string(),
            hrx::loom::Specialization::new("gemm"),
        )
    } else {
        (
            std::fs::read_to_string(&args[1])?,
            hrx::loom::Specialization::new(&args[2])
                .with_config("gemm.m", m.to_string())
                .with_config("gemm.k", k.to_string())
                .with_config("gemm.n", n.to_string()),
        )
    };
    let artifact = hrx::loom::Compiler::for_target(None, device.target())?
        .module(&source)
        .compile(&spec)?;
    let npu = unsafe {
        device.load_artifact(
            &artifact,
            1,
            KernelContract {
                bindings: vec![
                    binding(a_bytes, Access::Read),
                    binding(b_bytes, Access::Read),
                    binding(c_bytes, Access::Write),
                ],
                constants: vec![],
            },
        )
    }?;
    let prep = gpu_kernel(
        &runtime,
        include_str!("../tests/kernels/euler.loom"),
        "krea2_euler",
        &[
            ("krea2.euler.grid_x", (m * k).div_ceil(256).to_string()),
            ("krea2.euler.grid_y", "1".into()),
        ],
        m * k,
        vec![
            binding(a_bytes, Access::ReadWrite),
            binding(a_bytes, Access::Read),
        ],
    )?;
    let epilogue = gpu_kernel(
        &runtime,
        include_str!("../tests/kernels/add_f32.loom"),
        "add_f32",
        &[("add.grid", (m * n).div_ceil(256).to_string())],
        m * n,
        vec![binding(c_bytes, Access::ReadWrite)],
    )?;
    let allocate = |bytes| runtime.allocate(bytes, MemoryPlacement::Shared(device.clone()));
    let weights = allocate(b_bytes)?;
    let ones = allocate(a_bytes)?;
    for buffer in [&weights, &ones] {
        for word in buffer.map_write()?.as_chunks_mut::<2>().0 {
            word.copy_from_slice(&0x3f80u16.to_le_bytes());
        }
    }
    let prepare = || -> Result<(ExecutableGraph, Buffer)> {
        let a = allocate(a_bytes)?;
        let c = allocate(c_bytes)?;
        let mut graph = runtime.graph();
        graph.fill(a.view(), 0)?;
        graph.gpu(&prep, &[a.view(), ones.view()])?;
        graph.npu(&npu, &[a.view(), weights.view(), c.view()])?;
        graph.gpu(&epilogue, &[c.view()])?;
        Ok((graph.prepare()?, c))
    };
    let (first, c0) = prepare()?;
    let (second, c1) = prepare()?;
    for _ in 0..5 {
        first.submit()?.wait()?;
        second.submit()?.wait()?;
    }
    let before = runtime.statistics();
    let mut latency = Vec::with_capacity(50);
    for _ in 0..50 {
        let start = Instant::now();
        first.submit()?.wait()?;
        latency.push(start.elapsed().as_secs_f64());
    }
    let start = Instant::now();
    for _ in 0..50 {
        let a = first.submit()?;
        let b = second.submit()?;
        a.wait()?;
        b.wait()?;
    }
    let requests_per_second = 100.0 / start.elapsed().as_secs_f64();
    for c in [&c0, &c1] {
        if c.map_read()?
            .as_chunks::<4>()
            .0
            .iter()
            .any(|v| f32::from_le_bytes(*v) != k as f32 + 1.0)
        {
            return Err(hrx::Error::Message("GEMM pipeline output mismatch".into()));
        }
    }
    let after = runtime.statistics();
    assert_eq!(before.allocations, after.allocations);
    assert_eq!(before.imports, after.imports);
    assert_eq!(after.copied_bytes, 0);
    latency.sort_by(f64::total_cmp);
    println!(
        "M={m} K={k} N={n} p50_ms={:.3} p95_ms={:.3} pipelined_requests_s={requests_per_second:.2}",
        percentile(&latency, 50) * 1e3,
        percentile(&latency, 95) * 1e3
    );
    println!("{}", serde_json::to_string(&after)?);
    Ok(())
}

examples/gpu_npu_parallel.rs (line 199)

fn main() -> Result<()> {
    let args: Vec<_> = std::env::args().skip(1).collect();
    if args.len() > 4 || args.iter().any(|arg| arg == "--help") {
        println!(
            "usage: gpu_npu_parallel [gpu-elements=16777216] [npu-matrices=1024] [samples=21] [trace.json]"
        );
        return if args.len() <= 4 {
            Ok(())
        } else {
            Err(Error::Message("too many arguments".into()))
        };
    }
    let number = |index: usize, default: usize| -> Result<usize> {
        args.get(index).map_or(Ok(default), |value| {
            value
                .parse()
                .map_err(|_| Error::Message(format!("invalid integer: {value}")))
        })
    };
    let elements = number(0, 1 << 24)?;
    let matrices = number(1, 1024)?;
    let samples = number(2, 21)?;
    if !(256..=1 << 26).contains(&elements)
        || !elements.is_multiple_of(256)
        || !(1..=16384).contains(&matrices)
        || !(3..=1000).contains(&samples)
    {
        return Err(Error::Message("gpu-elements must be a multiple of 256 in 256..=67108864; npu-matrices in 1..=16384; samples in 3..=1000".into()));
    }
    let trace_path = args.get(3).map(PathBuf::from);
    let runtime = Runtime::new()?;
    let gpu = runtime.gpu()?;
    let npu = runtime.npu(0)?;
    println!(
        "GPU: {}; NPU: {}",
        gpu.target().as_str(),
        npu.target().as_str()
    );
    let gpu_compiler = Compiler::for_target(None, gpu.target())?;
    let gpu_artifact = gpu_compiler
        .import_cxx(CxxSource::new(
            "affine.cpp",
            format!(
                r#"
#include <hip/hip_runtime.h>
__global__ [[loom::workgroup_size(256, 1, 1), loom::workgroup_count({}, 1, 1)]]
void affine(const int* input, int* output) {{
  unsigned i = blockIdx.x * 256u + threadIdx.x;
  output[i] = input[i] * 3 + 7;
}}
"#,
                elements / 256
            ),
        ))?
        .compile(&Specialization::new("affine"))?;
    // SAFETY: the fixed launch covers exactly elements i32 values in each
    // nonaliasing binding; bounded input values cannot overflow the arithmetic.
    let gpu_kernel = unsafe {
        runtime.load_gpu_artifact(
            &gpu_artifact,
            &[],
            contract(&[(elements * 4, Access::Read), (elements * 4, Access::Write)]),
        )
    }?;
    let npu_artifact = Compiler::for_target(None, npu.target())?
        .module(include_str!("kernels/batched_gemm.xdna.loom"))
        .compile(&Specialization::new("gemm").with_config("gemm.batch", matrices.to_string()))?;
    // SAFETY: the pipeline consumes matrices pairs of 8x8 BF16 inputs and
    // produces the same number of 8x8 FP32 outputs on one NPU column.
    let npu_kernel = unsafe {
        npu.load_artifact(
            &npu_artifact,
            1,
            contract(&[
                (matrices * 128, Access::Read),
                (matrices * 128, Access::Read),
                (matrices * 256, Access::Write),
            ]),
        )
    }?;
    let input = runtime.allocate(elements * 4, MemoryPlacement::GpuLocal)?;
    let output = runtime.allocate(elements * 4, MemoryPlacement::GpuLocal)?;
    let host_input = runtime.allocate(elements * 4, MemoryPlacement::HostVisible)?;
    let host_output = runtime.allocate(elements * 4, MemoryPlacement::HostVisible)?;
    let mut upload = runtime.graph();
    upload.copy(input.view(), host_input.view())?;
    let upload = upload.prepare()?;
    let mut download = runtime.graph();
    download.copy(host_output.view(), output.view())?;
    let download = download.prepare()?;
    let a = runtime.allocate(matrices * 128, MemoryPlacement::NpuLocal(npu.clone()))?;
    let b = runtime.allocate(matrices * 128, MemoryPlacement::NpuLocal(npu.clone()))?;
    let c = runtime.allocate(matrices * 256, MemoryPlacement::NpuLocal(npu))?;

    let mut gpu_only = runtime.graph();
    gpu_only.gpu(&gpu_kernel, &[input.view(), output.view()])?;
    let gpu_only = gpu_only.prepare()?;
    let mut npu_only = runtime.graph();
    npu_only.npu(&npu_kernel, &[a.view(), b.view(), c.view()])?;
    let npu_only = npu_only.prepare()?;

    // Disjoint buffers imply no dependency edge. One submit admits both
    // branches to separate GPU/NPU workers; one wait joins their completions.
    let mut parallel = runtime.graph();
    parallel.gpu(&gpu_kernel, &[input.view(), output.view()])?;
    parallel.npu(&npu_kernel, &[a.view(), b.view(), c.view()])?;
    let parallel = parallel.prepare()?;
    for pass in 0..3 {
        initialize(&host_input, &a, &b, pass * 2)?;
        upload.submit()?.wait()?;
        gpu_only.submit()?.wait()?;
        npu_only.submit()?.wait()?;
        download.submit()?.wait()?;
        validate(&host_output, &c, pass * 2)?;
        // Change both inputs so validation cannot accept the sequential run's
        // old output if either parallel branch fails to write its result.
        initialize(&host_input, &a, &b, pass * 2 + 1)?;
        upload.submit()?.wait()?;
        parallel.submit()?.wait()?;
        download.submit()?.wait()?;
        validate(&host_output, &c, pass * 2 + 1)?;
    }

    let before = runtime.statistics();
    let mut sequential_ms = Vec::with_capacity(samples);
    let mut parallel_ms = Vec::with_capacity(samples);
    for sample in 0..samples {
        // Alternate order to reduce systematic clock/thermal ordering bias.
        for concurrent in if sample % 2 == 0 {
            [false, true]
        } else {
            [true, false]
        } {
            if concurrent {
                parallel_ms.push(completed_ms(&parallel)?);
            } else {
                let start = Instant::now();
                gpu_only.submit()?.wait()?;
                npu_only.submit()?.wait()?;
                sequential_ms.push(start.elapsed().as_secs_f64() * 1000.0);
            }
        }
    }
    let after = runtime.statistics();
    if before.allocations != after.allocations || before.imports != after.imports {
        return Err(Error::Message(
            "replay unexpectedly allocated or imported storage".into(),
        ));
    }
    download.submit()?.wait()?;
    validate(&host_output, &c, 5)?;
    let sequential = median(&mut sequential_ms);
    let concurrent = median(&mut parallel_ms);
    println!("{elements} GPU elements + {matrices} NPU matrices; {samples} paired samples");
    println!(
        "sequential median: {sequential:.3} ms; parallel median: {concurrent:.3} ms; speedup: {:.2}x",
        sequential / concurrent
    );
    println!(
        "Both CPU references match exactly, including changed-input replays. No native allocations/imports during timed replay."
    );

    // Capture separately so tracing does not affect the benchmark samples.
    runtime.start_trace(32)?;
    for _ in 0..4 {
        parallel.submit()?.wait()?;
    }
    let trace = runtime
        .finish_trace()
        .ok_or_else(|| Error::Message("missing execution trace".into()))?;
    let overlap: f64 = trace
        .events
        .iter()
        .filter(|e| e.lane == 1)
        .flat_map(|gpu| {
            trace.events.iter().filter(|e| e.lane == 3).map(move |npu| {
                ((gpu.start_us + gpu.duration_us).min(npu.start_us + npu.duration_us)
                    - gpu.start_us.max(npu.start_us))
                .max(0.0)
            })
        })
        .sum();
    println!(
        "GPU/NPU host-observed overlap across four replays: {overlap:.1} us (not device timestamps)"
    );
    if let Some(path) = trace_path {
        std::fs::write(&path, trace.to_json()?)?;
        println!("Trace: {}", path.display());
    }
    Ok(())
}

pub unsafe fn gpu_aliasing( &mut self, kernel: &GpuKernel, bindings: &[BufferView], ) -> Result<Node>

Invoke a trusted GPU kernel permitting writable argument aliases. Extents, alignment, runtime identity and cross-operation hazards remain checked.

Safety

The kernel must be valid for the overlapping argument regions supplied; no inter-thread data races or unsupported in-place accesses may occur.

pub fn npu( &mut self, kernel: &NpuKernel, bindings: &[BufferView], ) -> Result<Node>

Invoke a trusted NPU specialization with checked device and memory-bank compatibility. Host-only and imported BOs may originate in a different program context on the same device and memory bank. Their owning contexts remain retained.

Examples found in repository

examples/shared_roundtrip.rs (line 50)

fn main() -> Result<()> {
    let args: Vec<_> = std::env::args().collect();
    if args.len() > 2 {
        return Err(hrx::Error::Message(
            "usage: shared_roundtrip [bytes]".into(),
        ));
    }
    let bytes = args
        .get(1)
        .map(String::as_str)
        .unwrap_or("1048576")
        .parse::<usize>()
        .map_err(|e| hrx::Error::Message(e.to_string()))?;
    if bytes == 0 || bytes % 4096 != 0 || bytes > 64 * 1024 * 1024 {
        return Err(hrx::Error::Message(
            "bytes must be a multiple of 4096 in 4096..=67108864".into(),
        ));
    }
    let runtime = Runtime::new()?;
    let device = runtime.npu(0)?;
    let artifact = hrx::loom::Compiler::for_target(None, device.target())?
        .module(include_str!("../tests/kernels/copy.xdna.loom"))
        .compile(
            &hrx::loom::Specialization::new("copy")
                .with_config("copy.packets", (bytes / 4096).to_string()),
        )?;
    let binding = |bytes, access| BindingContract {
        bytes,
        access,
        alignment: 4,
        layout: "u32 passthrough".into(),
    };
    let contract = KernelContract {
        bindings: vec![binding(bytes, Access::Read), binding(bytes, Access::Write)],
        constants: vec![],
    };
    let kernel = unsafe { device.load_artifact(&artifact, 1, contract) }?;
    let a = runtime.allocate(bytes, MemoryPlacement::Shared(device.clone()))?;
    let c = runtime.allocate(bytes, MemoryPlacement::Shared(device.clone()))?;
    let result = runtime.allocate(bytes, MemoryPlacement::Shared(device.clone()))?;
    let mut graph = runtime.graph();
    graph.fill(a.view(), 0x5a)?;
    graph.npu(&kernel, &[a.view(), c.view()])?;
    graph.copy(result.view(), c.view())?;
    let graph = graph.prepare()?;
    for _ in 0..10 {
        graph.submit()?.wait()?;
        let mapped = result.map_read()?;
        if mapped.iter().any(|&b| b != 0x5a) {
            return Err(hrx::Error::Message("GPU/NPU roundtrip mismatch".into()));
        }
    }
    println!(
        "GPU -> NPU -> GPU: all {bytes} bytes correct, 10 executions, {} prepared regions",
        graph.regions()
    );
    Ok(())
}

examples/shared_bench.rs (line 53)

fn main() -> Result<()> {
    let args: Vec<_> = std::env::args().collect();
    if args.len() > 2 {
        return Err(hrx::Error::Message("usage: shared_bench [bytes]".into()));
    }
    let bytes = args
        .get(1)
        .map(String::as_str)
        .unwrap_or("1048576")
        .parse::<usize>()
        .map_err(|e| hrx::Error::Message(e.to_string()))?;
    if bytes == 0 || bytes % 4096 != 0 || bytes > 64 * 1024 * 1024 {
        return Err(hrx::Error::Message(
            "bytes must be a multiple of 4096 in 4096..=67108864".into(),
        ));
    }
    let runtime = Runtime::new()?;
    let device = runtime.npu(0)?;
    let artifact = hrx::loom::Compiler::for_target(None, device.target())?
        .module(include_str!("../tests/kernels/copy.xdna.loom"))
        .compile(
            &hrx::loom::Specialization::new("copy")
                .with_config("copy.packets", (bytes / 4096).to_string()),
        )?;
    let binding = |bytes, access| BindingContract {
        bytes,
        access,
        alignment: 4,
        layout: "passthrough".into(),
    };
    let kernel = unsafe {
        device.load_artifact(
            &artifact,
            1,
            KernelContract {
                bindings: vec![binding(bytes, Access::Read), binding(bytes, Access::Write)],
                constants: vec![],
            },
        )
    }?;
    let a = runtime.allocate(bytes, MemoryPlacement::Shared(device.clone()))?;
    let c = runtime.allocate(bytes, MemoryPlacement::Shared(device.clone()))?;
    let mut graph = runtime.graph();
    graph.fill(a.view(), 0x35)?;
    graph.npu(&kernel, &[a.view(), c.view()])?;
    let graph = graph.prepare()?;
    for _ in 0..5 {
        graph.submit()?.wait()?;
    }
    let mut samples = Vec::new();
    for _ in 0..100 {
        let start = Instant::now();
        graph.submit()?.wait()?;
        samples.push(start.elapsed().as_secs_f64());
    }
    if c.map_read()?.iter().any(|&b| b != 0x35) {
        return Err(hrx::Error::Message("incorrect benchmark output".into()));
    }
    samples.sort_by(f64::total_cmp);
    println!(
        "bytes={bytes} p50_us={:.2} p95_us={:.2}",
        percentile(&samples, 50) * 1e6,
        percentile(&samples, 95) * 1e6
    );
    Ok(())
}

examples/gemm_pipeline.rs (line 161)

fn main() -> Result<()> {
    let args: Vec<_> = std::env::args().collect();
    let args = if args.len() == 1 {
        vec![
            args[0].clone(),
            String::new(),
            "gemm".into(),
            "8".into(),
            "8".into(),
            "8".into(),
        ]
    } else {
        args
    };
    if args.len() != 6 {
        return Err(hrx::Error::Message(
            "usage: gemm_pipeline <gemm.loom> <symbol> M K N".into(),
        ));
    }
    let parse = |i: usize| {
        args[i]
            .parse::<usize>()
            .map_err(|e| hrx::Error::Message(e.to_string()))
    };
    let (m, k, n) = (parse(3)?, parse(4)?, parse(5)?);
    let extent = |a: usize, b: usize, width: usize| {
        a.checked_mul(b)
            .filter(|&count| count > 0 && count <= (1 << 30))
            .and_then(|count| count.checked_mul(width))
            .ok_or_else(|| hrx::Error::Message("invalid matrix extent".into()))
    };
    let (a_bytes, b_bytes, c_bytes) = (extent(m, k, 2)?, extent(k, n, 2)?, extent(m, n, 4)?);
    let runtime = Runtime::new()?;
    let device = runtime.npu(0)?;
    let (source, spec) = if args[1].is_empty() {
        (
            include_str!("../tests/kernels/gemm_bf16.xdna.loom").to_string(),
            hrx::loom::Specialization::new("gemm"),
        )
    } else {
        (
            std::fs::read_to_string(&args[1])?,
            hrx::loom::Specialization::new(&args[2])
                .with_config("gemm.m", m.to_string())
                .with_config("gemm.k", k.to_string())
                .with_config("gemm.n", n.to_string()),
        )
    };
    let artifact = hrx::loom::Compiler::for_target(None, device.target())?
        .module(&source)
        .compile(&spec)?;
    let npu = unsafe {
        device.load_artifact(
            &artifact,
            1,
            KernelContract {
                bindings: vec![
                    binding(a_bytes, Access::Read),
                    binding(b_bytes, Access::Read),
                    binding(c_bytes, Access::Write),
                ],
                constants: vec![],
            },
        )
    }?;
    let prep = gpu_kernel(
        &runtime,
        include_str!("../tests/kernels/euler.loom"),
        "krea2_euler",
        &[
            ("krea2.euler.grid_x", (m * k).div_ceil(256).to_string()),
            ("krea2.euler.grid_y", "1".into()),
        ],
        m * k,
        vec![
            binding(a_bytes, Access::ReadWrite),
            binding(a_bytes, Access::Read),
        ],
    )?;
    let epilogue = gpu_kernel(
        &runtime,
        include_str!("../tests/kernels/add_f32.loom"),
        "add_f32",
        &[("add.grid", (m * n).div_ceil(256).to_string())],
        m * n,
        vec![binding(c_bytes, Access::ReadWrite)],
    )?;
    let allocate = |bytes| runtime.allocate(bytes, MemoryPlacement::Shared(device.clone()));
    let weights = allocate(b_bytes)?;
    let ones = allocate(a_bytes)?;
    for buffer in [&weights, &ones] {
        for word in buffer.map_write()?.as_chunks_mut::<2>().0 {
            word.copy_from_slice(&0x3f80u16.to_le_bytes());
        }
    }
    let prepare = || -> Result<(ExecutableGraph, Buffer)> {
        let a = allocate(a_bytes)?;
        let c = allocate(c_bytes)?;
        let mut graph = runtime.graph();
        graph.fill(a.view(), 0)?;
        graph.gpu(&prep, &[a.view(), ones.view()])?;
        graph.npu(&npu, &[a.view(), weights.view(), c.view()])?;
        graph.gpu(&epilogue, &[c.view()])?;
        Ok((graph.prepare()?, c))
    };
    let (first, c0) = prepare()?;
    let (second, c1) = prepare()?;
    for _ in 0..5 {
        first.submit()?.wait()?;
        second.submit()?.wait()?;
    }
    let before = runtime.statistics();
    let mut latency = Vec::with_capacity(50);
    for _ in 0..50 {
        let start = Instant::now();
        first.submit()?.wait()?;
        latency.push(start.elapsed().as_secs_f64());
    }
    let start = Instant::now();
    for _ in 0..50 {
        let a = first.submit()?;
        let b = second.submit()?;
        a.wait()?;
        b.wait()?;
    }
    let requests_per_second = 100.0 / start.elapsed().as_secs_f64();
    for c in [&c0, &c1] {
        if c.map_read()?
            .as_chunks::<4>()
            .0
            .iter()
            .any(|v| f32::from_le_bytes(*v) != k as f32 + 1.0)
        {
            return Err(hrx::Error::Message("GEMM pipeline output mismatch".into()));
        }
    }
    let after = runtime.statistics();
    assert_eq!(before.allocations, after.allocations);
    assert_eq!(before.imports, after.imports);
    assert_eq!(after.copied_bytes, 0);
    latency.sort_by(f64::total_cmp);
    println!(
        "M={m} K={k} N={n} p50_ms={:.3} p95_ms={:.3} pipelined_requests_s={requests_per_second:.2}",
        percentile(&latency, 50) * 1e3,
        percentile(&latency, 95) * 1e3
    );
    println!("{}", serde_json::to_string(&after)?);
    Ok(())
}

examples/gpu_npu_parallel.rs (line 202)

fn main() -> Result<()> {
    let args: Vec<_> = std::env::args().skip(1).collect();
    if args.len() > 4 || args.iter().any(|arg| arg == "--help") {
        println!(
            "usage: gpu_npu_parallel [gpu-elements=16777216] [npu-matrices=1024] [samples=21] [trace.json]"
        );
        return if args.len() <= 4 {
            Ok(())
        } else {
            Err(Error::Message("too many arguments".into()))
        };
    }
    let number = |index: usize, default: usize| -> Result<usize> {
        args.get(index).map_or(Ok(default), |value| {
            value
                .parse()
                .map_err(|_| Error::Message(format!("invalid integer: {value}")))
        })
    };
    let elements = number(0, 1 << 24)?;
    let matrices = number(1, 1024)?;
    let samples = number(2, 21)?;
    if !(256..=1 << 26).contains(&elements)
        || !elements.is_multiple_of(256)
        || !(1..=16384).contains(&matrices)
        || !(3..=1000).contains(&samples)
    {
        return Err(Error::Message("gpu-elements must be a multiple of 256 in 256..=67108864; npu-matrices in 1..=16384; samples in 3..=1000".into()));
    }
    let trace_path = args.get(3).map(PathBuf::from);
    let runtime = Runtime::new()?;
    let gpu = runtime.gpu()?;
    let npu = runtime.npu(0)?;
    println!(
        "GPU: {}; NPU: {}",
        gpu.target().as_str(),
        npu.target().as_str()
    );
    let gpu_compiler = Compiler::for_target(None, gpu.target())?;
    let gpu_artifact = gpu_compiler
        .import_cxx(CxxSource::new(
            "affine.cpp",
            format!(
                r#"
#include <hip/hip_runtime.h>
__global__ [[loom::workgroup_size(256, 1, 1), loom::workgroup_count({}, 1, 1)]]
void affine(const int* input, int* output) {{
  unsigned i = blockIdx.x * 256u + threadIdx.x;
  output[i] = input[i] * 3 + 7;
}}
"#,
                elements / 256
            ),
        ))?
        .compile(&Specialization::new("affine"))?;
    // SAFETY: the fixed launch covers exactly elements i32 values in each
    // nonaliasing binding; bounded input values cannot overflow the arithmetic.
    let gpu_kernel = unsafe {
        runtime.load_gpu_artifact(
            &gpu_artifact,
            &[],
            contract(&[(elements * 4, Access::Read), (elements * 4, Access::Write)]),
        )
    }?;
    let npu_artifact = Compiler::for_target(None, npu.target())?
        .module(include_str!("kernels/batched_gemm.xdna.loom"))
        .compile(&Specialization::new("gemm").with_config("gemm.batch", matrices.to_string()))?;
    // SAFETY: the pipeline consumes matrices pairs of 8x8 BF16 inputs and
    // produces the same number of 8x8 FP32 outputs on one NPU column.
    let npu_kernel = unsafe {
        npu.load_artifact(
            &npu_artifact,
            1,
            contract(&[
                (matrices * 128, Access::Read),
                (matrices * 128, Access::Read),
                (matrices * 256, Access::Write),
            ]),
        )
    }?;
    let input = runtime.allocate(elements * 4, MemoryPlacement::GpuLocal)?;
    let output = runtime.allocate(elements * 4, MemoryPlacement::GpuLocal)?;
    let host_input = runtime.allocate(elements * 4, MemoryPlacement::HostVisible)?;
    let host_output = runtime.allocate(elements * 4, MemoryPlacement::HostVisible)?;
    let mut upload = runtime.graph();
    upload.copy(input.view(), host_input.view())?;
    let upload = upload.prepare()?;
    let mut download = runtime.graph();
    download.copy(host_output.view(), output.view())?;
    let download = download.prepare()?;
    let a = runtime.allocate(matrices * 128, MemoryPlacement::NpuLocal(npu.clone()))?;
    let b = runtime.allocate(matrices * 128, MemoryPlacement::NpuLocal(npu.clone()))?;
    let c = runtime.allocate(matrices * 256, MemoryPlacement::NpuLocal(npu))?;

    let mut gpu_only = runtime.graph();
    gpu_only.gpu(&gpu_kernel, &[input.view(), output.view()])?;
    let gpu_only = gpu_only.prepare()?;
    let mut npu_only = runtime.graph();
    npu_only.npu(&npu_kernel, &[a.view(), b.view(), c.view()])?;
    let npu_only = npu_only.prepare()?;

    // Disjoint buffers imply no dependency edge. One submit admits both
    // branches to separate GPU/NPU workers; one wait joins their completions.
    let mut parallel = runtime.graph();
    parallel.gpu(&gpu_kernel, &[input.view(), output.view()])?;
    parallel.npu(&npu_kernel, &[a.view(), b.view(), c.view()])?;
    let parallel = parallel.prepare()?;
    for pass in 0..3 {
        initialize(&host_input, &a, &b, pass * 2)?;
        upload.submit()?.wait()?;
        gpu_only.submit()?.wait()?;
        npu_only.submit()?.wait()?;
        download.submit()?.wait()?;
        validate(&host_output, &c, pass * 2)?;
        // Change both inputs so validation cannot accept the sequential run's
        // old output if either parallel branch fails to write its result.
        initialize(&host_input, &a, &b, pass * 2 + 1)?;
        upload.submit()?.wait()?;
        parallel.submit()?.wait()?;
        download.submit()?.wait()?;
        validate(&host_output, &c, pass * 2 + 1)?;
    }

    let before = runtime.statistics();
    let mut sequential_ms = Vec::with_capacity(samples);
    let mut parallel_ms = Vec::with_capacity(samples);
    for sample in 0..samples {
        // Alternate order to reduce systematic clock/thermal ordering bias.
        for concurrent in if sample % 2 == 0 {
            [false, true]
        } else {
            [true, false]
        } {
            if concurrent {
                parallel_ms.push(completed_ms(&parallel)?);
            } else {
                let start = Instant::now();
                gpu_only.submit()?.wait()?;
                npu_only.submit()?.wait()?;
                sequential_ms.push(start.elapsed().as_secs_f64() * 1000.0);
            }
        }
    }
    let after = runtime.statistics();
    if before.allocations != after.allocations || before.imports != after.imports {
        return Err(Error::Message(
            "replay unexpectedly allocated or imported storage".into(),
        ));
    }
    download.submit()?.wait()?;
    validate(&host_output, &c, 5)?;
    let sequential = median(&mut sequential_ms);
    let concurrent = median(&mut parallel_ms);
    println!("{elements} GPU elements + {matrices} NPU matrices; {samples} paired samples");
    println!(
        "sequential median: {sequential:.3} ms; parallel median: {concurrent:.3} ms; speedup: {:.2}x",
        sequential / concurrent
    );
    println!(
        "Both CPU references match exactly, including changed-input replays. No native allocations/imports during timed replay."
    );

    // Capture separately so tracing does not affect the benchmark samples.
    runtime.start_trace(32)?;
    for _ in 0..4 {
        parallel.submit()?.wait()?;
    }
    let trace = runtime
        .finish_trace()
        .ok_or_else(|| Error::Message("missing execution trace".into()))?;
    let overlap: f64 = trace
        .events
        .iter()
        .filter(|e| e.lane == 1)
        .flat_map(|gpu| {
            trace.events.iter().filter(|e| e.lane == 3).map(move |npu| {
                ((gpu.start_us + gpu.duration_us).min(npu.start_us + npu.duration_us)
                    - gpu.start_us.max(npu.start_us))
                .max(0.0)
            })
        })
        .sum();
    println!(
        "GPU/NPU host-observed overlap across four replays: {overlap:.1} us (not device timestamps)"
    );
    if let Some(path) = trace_path {
        std::fs::write(&path, trace.to_json()?)?;
        println!("Trace: {}", path.display());
    }
    Ok(())
}

pub fn prepare(self) -> Result<ExecutableGraph>

Lower GPU regions and import NPU views once; reserve reusable run slots.

Examples found in repository

examples/graph_prepare.rs (line 15)

fn main() -> hrx::Result<()> {
    let runtime = Runtime::new()?;
    let buffer = runtime.allocate(4096, MemoryPlacement::GpuLocal)?;
    let stream = Stream::open()?;
    let raw = stream.allocate(4096)?;
    for count in [128, 512, 1024, 2048] {
        let begin = Instant::now();
        let mut graph = runtime.graph();
        for _ in 0..count {
            graph.fill(buffer.view(), 0)?;
        }
        let graph = graph.prepare()?;
        let inferred = begin.elapsed();
        let begin = Instant::now();
        let mut native = stream.graph()?;
        let mut last = None;
        for _ in 0..count {
            last = Some(native.fill(last.as_slice(), raw.binding(), 0)?);
        }
        let _native = native.finish()?;
        println!(
            "nodes={count} inferred_us={} explicit_us={}",
            inferred.as_micros(),
            begin.elapsed().as_micros()
        );
        graph.submit()?.wait()?;
    }
    Ok(())
}

examples/shared_roundtrip.rs (line 52)

fn main() -> Result<()> {
    let args: Vec<_> = std::env::args().collect();
    if args.len() > 2 {
        return Err(hrx::Error::Message(
            "usage: shared_roundtrip [bytes]".into(),
        ));
    }
    let bytes = args
        .get(1)
        .map(String::as_str)
        .unwrap_or("1048576")
        .parse::<usize>()
        .map_err(|e| hrx::Error::Message(e.to_string()))?;
    if bytes == 0 || bytes % 4096 != 0 || bytes > 64 * 1024 * 1024 {
        return Err(hrx::Error::Message(
            "bytes must be a multiple of 4096 in 4096..=67108864".into(),
        ));
    }
    let runtime = Runtime::new()?;
    let device = runtime.npu(0)?;
    let artifact = hrx::loom::Compiler::for_target(None, device.target())?
        .module(include_str!("../tests/kernels/copy.xdna.loom"))
        .compile(
            &hrx::loom::Specialization::new("copy")
                .with_config("copy.packets", (bytes / 4096).to_string()),
        )?;
    let binding = |bytes, access| BindingContract {
        bytes,
        access,
        alignment: 4,
        layout: "u32 passthrough".into(),
    };
    let contract = KernelContract {
        bindings: vec![binding(bytes, Access::Read), binding(bytes, Access::Write)],
        constants: vec![],
    };
    let kernel = unsafe { device.load_artifact(&artifact, 1, contract) }?;
    let a = runtime.allocate(bytes, MemoryPlacement::Shared(device.clone()))?;
    let c = runtime.allocate(bytes, MemoryPlacement::Shared(device.clone()))?;
    let result = runtime.allocate(bytes, MemoryPlacement::Shared(device.clone()))?;
    let mut graph = runtime.graph();
    graph.fill(a.view(), 0x5a)?;
    graph.npu(&kernel, &[a.view(), c.view()])?;
    graph.copy(result.view(), c.view())?;
    let graph = graph.prepare()?;
    for _ in 0..10 {
        graph.submit()?.wait()?;
        let mapped = result.map_read()?;
        if mapped.iter().any(|&b| b != 0x5a) {
            return Err(hrx::Error::Message("GPU/NPU roundtrip mismatch".into()));
        }
    }
    println!(
        "GPU -> NPU -> GPU: all {bytes} bytes correct, 10 executions, {} prepared regions",
        graph.regions()
    );
    Ok(())
}

examples/shared_bench.rs (line 54)

fn main() -> Result<()> {
    let args: Vec<_> = std::env::args().collect();
    if args.len() > 2 {
        return Err(hrx::Error::Message("usage: shared_bench [bytes]".into()));
    }
    let bytes = args
        .get(1)
        .map(String::as_str)
        .unwrap_or("1048576")
        .parse::<usize>()
        .map_err(|e| hrx::Error::Message(e.to_string()))?;
    if bytes == 0 || bytes % 4096 != 0 || bytes > 64 * 1024 * 1024 {
        return Err(hrx::Error::Message(
            "bytes must be a multiple of 4096 in 4096..=67108864".into(),
        ));
    }
    let runtime = Runtime::new()?;
    let device = runtime.npu(0)?;
    let artifact = hrx::loom::Compiler::for_target(None, device.target())?
        .module(include_str!("../tests/kernels/copy.xdna.loom"))
        .compile(
            &hrx::loom::Specialization::new("copy")
                .with_config("copy.packets", (bytes / 4096).to_string()),
        )?;
    let binding = |bytes, access| BindingContract {
        bytes,
        access,
        alignment: 4,
        layout: "passthrough".into(),
    };
    let kernel = unsafe {
        device.load_artifact(
            &artifact,
            1,
            KernelContract {
                bindings: vec![binding(bytes, Access::Read), binding(bytes, Access::Write)],
                constants: vec![],
            },
        )
    }?;
    let a = runtime.allocate(bytes, MemoryPlacement::Shared(device.clone()))?;
    let c = runtime.allocate(bytes, MemoryPlacement::Shared(device.clone()))?;
    let mut graph = runtime.graph();
    graph.fill(a.view(), 0x35)?;
    graph.npu(&kernel, &[a.view(), c.view()])?;
    let graph = graph.prepare()?;
    for _ in 0..5 {
        graph.submit()?.wait()?;
    }
    let mut samples = Vec::new();
    for _ in 0..100 {
        let start = Instant::now();
        graph.submit()?.wait()?;
        samples.push(start.elapsed().as_secs_f64());
    }
    if c.map_read()?.iter().any(|&b| b != 0x35) {
        return Err(hrx::Error::Message("incorrect benchmark output".into()));
    }
    samples.sort_by(f64::total_cmp);
    println!(
        "bytes={bytes} p50_us={:.2} p95_us={:.2}",
        percentile(&samples, 50) * 1e6,
        percentile(&samples, 95) * 1e6
    );
    Ok(())
}

examples/gemm_pipeline.rs (line 163)

fn main() -> Result<()> {
    let args: Vec<_> = std::env::args().collect();
    let args = if args.len() == 1 {
        vec![
            args[0].clone(),
            String::new(),
            "gemm".into(),
            "8".into(),
            "8".into(),
            "8".into(),
        ]
    } else {
        args
    };
    if args.len() != 6 {
        return Err(hrx::Error::Message(
            "usage: gemm_pipeline <gemm.loom> <symbol> M K N".into(),
        ));
    }
    let parse = |i: usize| {
        args[i]
            .parse::<usize>()
            .map_err(|e| hrx::Error::Message(e.to_string()))
    };
    let (m, k, n) = (parse(3)?, parse(4)?, parse(5)?);
    let extent = |a: usize, b: usize, width: usize| {
        a.checked_mul(b)
            .filter(|&count| count > 0 && count <= (1 << 30))
            .and_then(|count| count.checked_mul(width))
            .ok_or_else(|| hrx::Error::Message("invalid matrix extent".into()))
    };
    let (a_bytes, b_bytes, c_bytes) = (extent(m, k, 2)?, extent(k, n, 2)?, extent(m, n, 4)?);
    let runtime = Runtime::new()?;
    let device = runtime.npu(0)?;
    let (source, spec) = if args[1].is_empty() {
        (
            include_str!("../tests/kernels/gemm_bf16.xdna.loom").to_string(),
            hrx::loom::Specialization::new("gemm"),
        )
    } else {
        (
            std::fs::read_to_string(&args[1])?,
            hrx::loom::Specialization::new(&args[2])
                .with_config("gemm.m", m.to_string())
                .with_config("gemm.k", k.to_string())
                .with_config("gemm.n", n.to_string()),
        )
    };
    let artifact = hrx::loom::Compiler::for_target(None, device.target())?
        .module(&source)
        .compile(&spec)?;
    let npu = unsafe {
        device.load_artifact(
            &artifact,
            1,
            KernelContract {
                bindings: vec![
                    binding(a_bytes, Access::Read),
                    binding(b_bytes, Access::Read),
                    binding(c_bytes, Access::Write),
                ],
                constants: vec![],
            },
        )
    }?;
    let prep = gpu_kernel(
        &runtime,
        include_str!("../tests/kernels/euler.loom"),
        "krea2_euler",
        &[
            ("krea2.euler.grid_x", (m * k).div_ceil(256).to_string()),
            ("krea2.euler.grid_y", "1".into()),
        ],
        m * k,
        vec![
            binding(a_bytes, Access::ReadWrite),
            binding(a_bytes, Access::Read),
        ],
    )?;
    let epilogue = gpu_kernel(
        &runtime,
        include_str!("../tests/kernels/add_f32.loom"),
        "add_f32",
        &[("add.grid", (m * n).div_ceil(256).to_string())],
        m * n,
        vec![binding(c_bytes, Access::ReadWrite)],
    )?;
    let allocate = |bytes| runtime.allocate(bytes, MemoryPlacement::Shared(device.clone()));
    let weights = allocate(b_bytes)?;
    let ones = allocate(a_bytes)?;
    for buffer in [&weights, &ones] {
        for word in buffer.map_write()?.as_chunks_mut::<2>().0 {
            word.copy_from_slice(&0x3f80u16.to_le_bytes());
        }
    }
    let prepare = || -> Result<(ExecutableGraph, Buffer)> {
        let a = allocate(a_bytes)?;
        let c = allocate(c_bytes)?;
        let mut graph = runtime.graph();
        graph.fill(a.view(), 0)?;
        graph.gpu(&prep, &[a.view(), ones.view()])?;
        graph.npu(&npu, &[a.view(), weights.view(), c.view()])?;
        graph.gpu(&epilogue, &[c.view()])?;
        Ok((graph.prepare()?, c))
    };
    let (first, c0) = prepare()?;
    let (second, c1) = prepare()?;
    for _ in 0..5 {
        first.submit()?.wait()?;
        second.submit()?.wait()?;
    }
    let before = runtime.statistics();
    let mut latency = Vec::with_capacity(50);
    for _ in 0..50 {
        let start = Instant::now();
        first.submit()?.wait()?;
        latency.push(start.elapsed().as_secs_f64());
    }
    let start = Instant::now();
    for _ in 0..50 {
        let a = first.submit()?;
        let b = second.submit()?;
        a.wait()?;
        b.wait()?;
    }
    let requests_per_second = 100.0 / start.elapsed().as_secs_f64();
    for c in [&c0, &c1] {
        if c.map_read()?
            .as_chunks::<4>()
            .0
            .iter()
            .any(|v| f32::from_le_bytes(*v) != k as f32 + 1.0)
        {
            return Err(hrx::Error::Message("GEMM pipeline output mismatch".into()));
        }
    }
    let after = runtime.statistics();
    assert_eq!(before.allocations, after.allocations);
    assert_eq!(before.imports, after.imports);
    assert_eq!(after.copied_bytes, 0);
    latency.sort_by(f64::total_cmp);
    println!(
        "M={m} K={k} N={n} p50_ms={:.3} p95_ms={:.3} pipelined_requests_s={requests_per_second:.2}",
        percentile(&latency, 50) * 1e3,
        percentile(&latency, 95) * 1e3
    );
    println!("{}", serde_json::to_string(&after)?);
    Ok(())
}

examples/gpu_npu_parallel.rs (line 190)

fn main() -> Result<()> {
    let args: Vec<_> = std::env::args().skip(1).collect();
    if args.len() > 4 || args.iter().any(|arg| arg == "--help") {
        println!(
            "usage: gpu_npu_parallel [gpu-elements=16777216] [npu-matrices=1024] [samples=21] [trace.json]"
        );
        return if args.len() <= 4 {
            Ok(())
        } else {
            Err(Error::Message("too many arguments".into()))
        };
    }
    let number = |index: usize, default: usize| -> Result<usize> {
        args.get(index).map_or(Ok(default), |value| {
            value
                .parse()
                .map_err(|_| Error::Message(format!("invalid integer: {value}")))
        })
    };
    let elements = number(0, 1 << 24)?;
    let matrices = number(1, 1024)?;
    let samples = number(2, 21)?;
    if !(256..=1 << 26).contains(&elements)
        || !elements.is_multiple_of(256)
        || !(1..=16384).contains(&matrices)
        || !(3..=1000).contains(&samples)
    {
        return Err(Error::Message("gpu-elements must be a multiple of 256 in 256..=67108864; npu-matrices in 1..=16384; samples in 3..=1000".into()));
    }
    let trace_path = args.get(3).map(PathBuf::from);
    let runtime = Runtime::new()?;
    let gpu = runtime.gpu()?;
    let npu = runtime.npu(0)?;
    println!(
        "GPU: {}; NPU: {}",
        gpu.target().as_str(),
        npu.target().as_str()
    );
    let gpu_compiler = Compiler::for_target(None, gpu.target())?;
    let gpu_artifact = gpu_compiler
        .import_cxx(CxxSource::new(
            "affine.cpp",
            format!(
                r#"
#include <hip/hip_runtime.h>
__global__ [[loom::workgroup_size(256, 1, 1), loom::workgroup_count({}, 1, 1)]]
void affine(const int* input, int* output) {{
  unsigned i = blockIdx.x * 256u + threadIdx.x;
  output[i] = input[i] * 3 + 7;
}}
"#,
                elements / 256
            ),
        ))?
        .compile(&Specialization::new("affine"))?;
    // SAFETY: the fixed launch covers exactly elements i32 values in each
    // nonaliasing binding; bounded input values cannot overflow the arithmetic.
    let gpu_kernel = unsafe {
        runtime.load_gpu_artifact(
            &gpu_artifact,
            &[],
            contract(&[(elements * 4, Access::Read), (elements * 4, Access::Write)]),
        )
    }?;
    let npu_artifact = Compiler::for_target(None, npu.target())?
        .module(include_str!("kernels/batched_gemm.xdna.loom"))
        .compile(&Specialization::new("gemm").with_config("gemm.batch", matrices.to_string()))?;
    // SAFETY: the pipeline consumes matrices pairs of 8x8 BF16 inputs and
    // produces the same number of 8x8 FP32 outputs on one NPU column.
    let npu_kernel = unsafe {
        npu.load_artifact(
            &npu_artifact,
            1,
            contract(&[
                (matrices * 128, Access::Read),
                (matrices * 128, Access::Read),
                (matrices * 256, Access::Write),
            ]),
        )
    }?;
    let input = runtime.allocate(elements * 4, MemoryPlacement::GpuLocal)?;
    let output = runtime.allocate(elements * 4, MemoryPlacement::GpuLocal)?;
    let host_input = runtime.allocate(elements * 4, MemoryPlacement::HostVisible)?;
    let host_output = runtime.allocate(elements * 4, MemoryPlacement::HostVisible)?;
    let mut upload = runtime.graph();
    upload.copy(input.view(), host_input.view())?;
    let upload = upload.prepare()?;
    let mut download = runtime.graph();
    download.copy(host_output.view(), output.view())?;
    let download = download.prepare()?;
    let a = runtime.allocate(matrices * 128, MemoryPlacement::NpuLocal(npu.clone()))?;
    let b = runtime.allocate(matrices * 128, MemoryPlacement::NpuLocal(npu.clone()))?;
    let c = runtime.allocate(matrices * 256, MemoryPlacement::NpuLocal(npu))?;

    let mut gpu_only = runtime.graph();
    gpu_only.gpu(&gpu_kernel, &[input.view(), output.view()])?;
    let gpu_only = gpu_only.prepare()?;
    let mut npu_only = runtime.graph();
    npu_only.npu(&npu_kernel, &[a.view(), b.view(), c.view()])?;
    let npu_only = npu_only.prepare()?;

    // Disjoint buffers imply no dependency edge. One submit admits both
    // branches to separate GPU/NPU workers; one wait joins their completions.
    let mut parallel = runtime.graph();
    parallel.gpu(&gpu_kernel, &[input.view(), output.view()])?;
    parallel.npu(&npu_kernel, &[a.view(), b.view(), c.view()])?;
    let parallel = parallel.prepare()?;
    for pass in 0..3 {
        initialize(&host_input, &a, &b, pass * 2)?;
        upload.submit()?.wait()?;
        gpu_only.submit()?.wait()?;
        npu_only.submit()?.wait()?;
        download.submit()?.wait()?;
        validate(&host_output, &c, pass * 2)?;
        // Change both inputs so validation cannot accept the sequential run's
        // old output if either parallel branch fails to write its result.
        initialize(&host_input, &a, &b, pass * 2 + 1)?;
        upload.submit()?.wait()?;
        parallel.submit()?.wait()?;
        download.submit()?.wait()?;
        validate(&host_output, &c, pass * 2 + 1)?;
    }

    let before = runtime.statistics();
    let mut sequential_ms = Vec::with_capacity(samples);
    let mut parallel_ms = Vec::with_capacity(samples);
    for sample in 0..samples {
        // Alternate order to reduce systematic clock/thermal ordering bias.
        for concurrent in if sample % 2 == 0 {
            [false, true]
        } else {
            [true, false]
        } {
            if concurrent {
                parallel_ms.push(completed_ms(&parallel)?);
            } else {
                let start = Instant::now();
                gpu_only.submit()?.wait()?;
                npu_only.submit()?.wait()?;
                sequential_ms.push(start.elapsed().as_secs_f64() * 1000.0);
            }
        }
    }
    let after = runtime.statistics();
    if before.allocations != after.allocations || before.imports != after.imports {
        return Err(Error::Message(
            "replay unexpectedly allocated or imported storage".into(),
        ));
    }
    download.submit()?.wait()?;
    validate(&host_output, &c, 5)?;
    let sequential = median(&mut sequential_ms);
    let concurrent = median(&mut parallel_ms);
    println!("{elements} GPU elements + {matrices} NPU matrices; {samples} paired samples");
    println!(
        "sequential median: {sequential:.3} ms; parallel median: {concurrent:.3} ms; speedup: {:.2}x",
        sequential / concurrent
    );
    println!(
        "Both CPU references match exactly, including changed-input replays. No native allocations/imports during timed replay."
    );

    // Capture separately so tracing does not affect the benchmark samples.
    runtime.start_trace(32)?;
    for _ in 0..4 {
        parallel.submit()?.wait()?;
    }
    let trace = runtime
        .finish_trace()
        .ok_or_else(|| Error::Message("missing execution trace".into()))?;
    let overlap: f64 = trace
        .events
        .iter()
        .filter(|e| e.lane == 1)
        .flat_map(|gpu| {
            trace.events.iter().filter(|e| e.lane == 3).map(move |npu| {
                ((gpu.start_us + gpu.duration_us).min(npu.start_us + npu.duration_us)
                    - gpu.start_us.max(npu.start_us))
                .max(0.0)
            })
        })
        .sum();
    println!(
        "GPU/NPU host-observed overlap across four replays: {overlap:.1} us (not device timestamps)"
    );
    if let Some(path) = trace_path {
        std::fs::write(&path, trace.to_json()?)?;
        println!("Trace: {}", path.display());
    }
    Ok(())
}