amdgpu-device-libs 0.2.0

Support library for the amdgpu target, linking to device-libs
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95

# amdgpu-device-libs [![docs.rs](https://docs.rs/amdgpu-device-libs/badge.svg)](https://docs.rs/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](https://github.com/ROCm/llvm-project/tree/amd-staging/amd/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:
```toml
# 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`:
```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:
```rust
// build.rs
fn main() {
    amdgpu_device_libs_build::build();
}
```

### Example

Minimal usage sample, see [`examples/println`](https://github.com/Flakebi/amdgpu-rs/tree/main/examples/println) for the full code.
```rust
#![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];
    }
}
```