Struct Target 

pub struct Target(/* private fields */);

An exact AMDGPU or XDNA deployment key.

This is the profile target: what a device reports and what the compiler emits for. It is deliberately narrower than the target a Loom source file names in amdgpu.target<...>, which accepts generic families such as gfx11-generic. The two are chosen independently — generic source compiles under a bare profile — so this type rejects generic names rather than silently accepting one a device can never report.

Implementations

impl Target

pub fn new(key: &str) -> Result<Self>

Validate an architecture key, for example gfx1100, gfx1151 or gfx942. Availability is determined by the selected native runtime and compiler. Keys must be bare architecture names, without feature suffixes.

pub fn xdna() -> Self

Exact Strix Halo NPU5 deployment profile.

Examples found in repository

examples/compile_npu.rs (line 9)

fn main() -> Result<()> {
    let packets = std::env::args().nth(1).unwrap_or_else(|| "256".into());
    let artifact = Compiler::for_target(None, &Target::xdna())?
        .module(include_str!("../tests/kernels/copy.xdna.loom"))
        .compile(
            &Specialization::new("copy")
                .with_config("copy.packets", packets)
                .with_report(ReportMode::Summary),
        )?;
    println!("{}", artifact.path().display());
    Ok(())
}

pub fn is_xdna(&self) -> bool

Whether this profile targets native XDNA execution.

pub fn as_str(&self) -> &str

Exact architecture key.

Examples found in repository

examples/queue_pool_bench.rs (line 173)

fn main() -> Result<()> {
    let args: Vec<_> = std::env::args().skip(1).collect();
    let parse = |value: &str| value.parse::<usize>().ok().filter(|&n| n > 0);
    let queues = args.first().and_then(|arg| parse(arg));
    let samples = args.get(1).map_or(Some(15), |arg| parse(arg));
    let (queues, samples) = match (queues, samples, args.len()) {
        (Some(queues), Some(samples), 1..=2) => (queues, samples),
        _ => {
            return Err(hrx::Error::Message(
                "usage: queue_pool_bench QUEUES [SAMPLES] (positive integers)".into(),
            ));
        }
    };
    let device = Device::open(0)?;
    device.set_stream_queue_count(queues)?;
    let start = Instant::now();
    let streams = (0..STREAMS)
        .map(|_| device.stream())
        .collect::<Result<Vec<_>>>()?;
    let creation_ms = milliseconds(start);
    // Keep these streams alive so every case uses the same already-created queues.
    let mut small = (0..STREAMS)
        .map(|_| Lane::new(&device, 4096))
        .collect::<Result<Vec<_>>>()?;
    let small_batch = 1024;
    let small_times = throughput(&mut small, samples, small_batch)?;
    let threaded_times = concurrent(&mut small, samples, small_batch)?;
    let mut idle_latency = Vec::new();
    for round in 0..samples + WARMUP {
        for lane in &mut small {
            let start = Instant::now();
            lane.enqueue()?;
            lane.stream.synchronize()?;
            if round >= WARMUP {
                idle_latency.push(milliseconds(start));
            }
        }
    }
    // A 1 GiB aggregate working set exceeds the on-device caches on gfx1151.
    let bulk_bytes = 32 * 1024 * 1024;
    let bulk_batch = 16;
    let mut bulk = (0..STREAMS)
        .map(|_| Lane::new(&device, bulk_bytes))
        .collect::<Result<Vec<_>>>()?;
    let bulk_times = throughput(&mut bulk, samples, bulk_batch)?;
    let mut loaded_latency = Vec::new();
    for round in 0..samples + WARMUP {
        // Probe every queue position. Enqueue a bulk backlog on all streams,
        // then measure one small copy's completion before draining the backlog.
        for probe in &mut small {
            for _ in 0..4 {
                for lane in &bulk {
                    lane.enqueue()?;
                }
            }
            let start = Instant::now();
            probe.enqueue()?;
            probe.stream.synchronize()?;
            if round >= WARMUP {
                loaded_latency.push(milliseconds(start));
            }
            for lane in &mut bulk {
                lane.stream.synchronize()?;
            }
        }
    }
    for lane in small.iter_mut().chain(bulk.iter_mut()) {
        lane.check()?;
    }
    let idle = Distribution::from_samples(idle_latency)?;
    let loaded = Distribution::from_samples(loaded_latency)?;
    println!(
        "{}",
        json!({
            "target": device.target().as_str(), "queues": queues, "streams": STREAMS,
            "samples": samples, "warmup": WARMUP, "stream_creation_ms": creation_ms,
            "small_commands_per_second": STREAMS as f64 * small_batch as f64 * 1_000.0 / small_times.median_ms,
            "threaded_commands_per_second": STREAMS as f64 * small_batch as f64 * 1_000.0 / threaded_times.median_ms,
            "bulk_copy_gb_s": 2.0 * bulk_bytes as f64 * STREAMS as f64 * bulk_batch as f64 / (bulk_times.median_ms * 1e6),
            "small_batch_median_ms": small_times.median_ms, "small_batch_p95_ms": small_times.p95_ms,
            "threaded_batch_median_ms": threaded_times.median_ms, "threaded_batch_p95_ms": threaded_times.p95_ms,
            "bulk_batch_median_ms": bulk_times.median_ms, "bulk_batch_p95_ms": bulk_times.p95_ms,
            "idle_latency_median_us": idle.median_ms * 1_000.0, "idle_latency_p95_us": idle.p95_ms * 1_000.0,
            "loaded_latency_median_us": loaded.median_ms * 1_000.0, "loaded_latency_p95_us": loaded.p95_ms * 1_000.0,
        })
    );
    drop(streams);
    Ok(())
}

examples/gpu_npu_parallel.rs (line 139)

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 manifest_key(&self) -> String

Bundle platform and architecture identifier.

Trait Implementations

impl Clone for Target

fn clone(&self) -> Self

Returns a duplicate of the value.

fn clone_from(&mut self, source: &Self)

Performs copy-assignment from source.

impl Debug for Target

fn fmt(&self, f: &mut Formatter<'_>) -> Result

Formats the value using the given formatter.

impl Default for Target

fn default() -> Self

Returns the “default value” for a type.

impl Eq for Target

impl PartialEq for Target

fn eq(&self, other: &Self) -> bool

Equality operator ==.

fn ne(&self, other: &Rhs) -> bool

Inequality operator !=.

impl StructuralPartialEq for Target