amdgpu-device-libs

Support library for the amdgpu target.

By default, the amdgpu target supports core, but not std. alloc is supported when a global allocator is specified.

amdgpu-device-libs brings some std-like features to the amdgpu target:

• print!() and println!() macros for printing on the host stdout
• A global allocator to support alloc
• A panic handler
• Access to more intrinsics and device-libs functions

All these features are enabled by default, but can be turned on selectively with default-features = false, features = […].

amdgpu-device-libs works by linking to the ROCm device-libs and a pre-compiled helper library. The libraries are linked from a ROCm installation. If libraries are found not found by default, set the environment variable HIP_DEVICE_LIB_PATH to the bitcode files. The path should end with amdgcn/bitcode and contain .bc files.

Usage

Create a new cargo library project and change it to compile a cdylib:

# Cargo.toml
# In old Rust versions (before #149624), lto needed to be forced with `[profile.dev] lto = true` and  `[profile.release] lto = true`
[lib]
# Compile a cdylib
crate-type = ["cdylib"]

[build-dependencies]
# Used in build script to specify linker flags and link in device-libs
amdgpu-device-libs-build = "<version>"

[dependencies]
amdgpu-device-libs = "<version>"

Add extra flags in .cargo/config.toml:

# .cargo/config.toml
[build]
target = "amdgcn-amd-amdhsa"
# Enable linker-plugin-lto
# Either add -Ctarget-cpu=gfx<version> here or specify it in CARGO_BUILD_RUSTFLAGS='-Ctarget-cpu=gfx<version>'
# Note that setting RUSTFLAGS overwrites the flags added here!
# In old Rust versions (before #136840), linker-plugin-lto is not working as expected and -Zemit-thin-lto=no needs to be added to rustflags
rustflags = ["-Clinker-plugin-lto"]

[unstable]
build-std = ["core", "alloc"]

And add a build.rs build script that links to the required libraries:

// build.rs
fn main() {
    amdgpu_device_libs_build::build();
}

Example

Minimal usage sample, see examples/println for the full code.

#![feature(abi_gpu_kernel)]
#![no_std]

extern crate alloc;

use alloc::vec::Vec;

use amdgpu_device_libs::prelude::*;

#[unsafe(no_mangle)]
pub extern "gpu-kernel" fn kernel(output: *mut u32) {
    let wg_id = workgroup_id_x();
    let id = workitem_id_x();
    let dispatch = dispatch_ptr();
    let complete_id = wg_id as usize * dispatch.workgroup_size_x as usize + id as usize;

    println!("Hello world from the GPU! (thread {wg_id}-{id})");

    let mut v = Vec::<u32>::new();
    for i in 0..100 {
        v.push(100 + i);
    }

    unsafe {
        *output.add(complete_id) = v[complete_id];
    }
}