1 unstable release
Uses new Rust 2024
new 0.1.0 | Mar 24, 2025 |
---|
#312 in Hardware support
32KB
275 lines
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!()
andprintln!()
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.
To make sure the libraries are found, set the environment variable ROCM_PATH
or ROCM_DEVICE_LIB_PATH
(higher priority if it is set).
It looks for amdgcn/bitcode/*.bc
files in this path.
Usage
Create a new cargo library project and change it to compile a cdylib:
# Cargo.toml
# Force lto
[profile.dev]
lto = true
[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 = { path = "../../amdgpu-device-libs-build" }
[dependencies]
amdgpu-device-libs = { path = "../../amdgpu-device-libs" }
Add extra flags in .cargo/config.toml
:
# .cargo/config.toml
[build]
target = "amdgcn-amd-amdhsa"
# Enable linker-plugin-lto and workarounds
# Either add -Ctarget-cpu=gfx<version> here or specify it in CARGO_BUILD_RUSTFLAGS='-Ctarget-cpu=gfx<version>'
rustflags = ["-Clinker-plugin-lto", "-Zemit-thin-lto=no"]
[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];
}
}