hanzo/rocm-rs
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

feat: update dependencies and versioning, add new examples and kernel attributes

Browse Source
This commit is contained in:
PTFOPlayer
2026-03-25 08:51:46 +01:00
parent 391ee4065a
commit 85f3c88faf
13 changed files with 253 additions and 72 deletions
Download Patch File Download Diff File Expand all files Collapse all files
.gitignore
+1
View File
@@ -1,5 +1,6 @@
/target
**/target
**/kernel_sources
**/__build_in_kernels_sorting
**/bindings.rs
*.code-workspace
Cargo.lock
Generated
+88 -3
View File
@@ -86,6 +86,27 @@ dependencies = [
"libloading",
]
[[package]]
name = "csv"
version = "1.4.0"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "52cd9d68cf7efc6ddfaaee42e7288d3a99d613d4b50f76ce9827ae0c6e14f938"
dependencies = [
"csv-core",
"itoa",
"ryu",
"serde_core",
]
[[package]]
name = "csv-core"
version = "0.1.13"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "704a3c26996a80471189265814dbc2c257598b96b8a7feae2d31ace646bb9782"
dependencies = [
"memchr",
]
[[package]]
name = "either"
version = "1.15.0"
@@ -102,6 +123,13 @@ dependencies = [
"windows-sys",
]
[[package]]
name = "example"
version = "0.1.0"
dependencies = [
"rocm-rs",
]
[[package]]
name = "fslock"
version = "0.2.1"
@@ -203,6 +231,21 @@ version = "0.2.1"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "68354c5c6bd36d73ff3feceb05efa59b6acb7626617f4962be322a825e61f79a"
[[package]]
name = "miopen_basic"
version = "0.1.0"
dependencies = [
"rocm-rs",
]
[[package]]
name = "multi_tensor"
version = "0.1.0"
dependencies = [
"csv",
"rocm-rs",
]
[[package]]
name = "nom"
version = "7.1.3"
@@ -282,9 +325,16 @@ version = "0.8.5"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "2b15c43186be67a4fd63bee50d0303afffcef381492ebe2c5d87f324e1b8815c"
[[package]]
name = "rocblas_basic"
version = "0.1.0"
dependencies = [
"rocm-rs",
]
[[package]]
name = "rocm-rs"
version = "0.5.0"
version = "0.5.1"
dependencies = [
"bindgen 0.71.1",
"paste",
@@ -294,9 +344,9 @@ dependencies = [
[[package]]
name = "rocm_kernel_macros"
version = "0.4.2"
version = "0.5.0"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "c180cfaeffccc518db1100beffb3eae2827ae52ad89c56a5417aa0df37f13ca3"
checksum = "2554ba75ab4ff9e6dd58ff73c746925baf5411e5642a9784efcbb8aa1b6d64bc"
dependencies = [
"fslock",
"proc-macro2",
@@ -327,6 +377,20 @@ dependencies = [
"libc",
]
[[package]]
name = "rust_kernel"
version = "0.1.0"
dependencies = [
"rocm-rs",
]
[[package]]
name = "rust_kernel_async"
version = "0.1.0"
dependencies = [
"rocm-rs",
]
[[package]]
name = "rustc-hash"
version = "1.1.0"
@@ -358,6 +422,13 @@ version = "1.0.20"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "28d3b2b1366ec20994f1fd18c3c594f05c5dd4bc44d8bb0c1c632c8d6829481f"
[[package]]
name = "saxpy"
version = "0.1.0"
dependencies = [
"rocm-rs",
]
[[package]]
name = "serde"
version = "1.0.226"
@@ -407,6 +478,13 @@ version = "1.3.0"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "0fda2ff0d084019ba4d7c6f371c95d8fd75ce3524c3cb8fb653a3023f6323e64"
[[package]]
name = "sort"
version = "0.1.0"
dependencies = [
"rocm-rs",
]
[[package]]
name = "syn"
version = "2.0.100"
@@ -424,6 +502,13 @@ version = "1.0.18"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "5a5f39404a5da50712a4c1eecf25e90dd62b613502b7e925fd4e4d19b5c96512"
[[package]]
name = "vector_add_example"
version = "0.1.0"
dependencies = [
"rocm-rs",
]
[[package]]
name = "which"
version = "4.4.2"
Cargo.toml
+16 -2
View File
@@ -1,6 +1,6 @@
[package]
name = "rocm-rs"
version = "0.5.0"
version = "0.5.1"
edition = "2024"
description = "Rust bindings for AMD ROCm libraries"
license = "MIT"
@@ -16,7 +16,7 @@ doctest = false
[dependencies]
rocm_smi_lib = { version = "0.3.1", optional = true }
rocm_kernel_macros = {version = "0.4.2", optional = true}
rocm_kernel_macros = {version = "0.5.0", optional = true}
paste = "1.0.15"
[build-dependencies]
@@ -29,3 +29,17 @@ miopen = []
rocprofiler = []
rocsolver = []
macros=["dep:rocm_kernel_macros"]
[workspace]
members = [
"src/hip/examples/rust_kernel",
"src/hip/examples/rust_kernel_async",
"src/hip/examples/saxpy",
"src/hip/examples/sort",
"src/hip/examples/vector_add",
"src/miopen/examples/basic",
"src/miopen/examples/multi_tensor",
"src/rocblas/examples/basic",
"src/rocrand/examples/normal",
]
exclude = ["__build_in_kernels_sorting", "kernel_sources"]
README.md
+88 -12
View File
@@ -38,7 +38,7 @@ Add this to your `Cargo.toml`:
```toml
[dependencies]
rocm-rs = "4.2"
rocm-rs = "5.1"
```
## Usage
@@ -50,8 +50,11 @@ First, ensure that the ROCm libraries are in your library path or set the `ROCM_
```rust
use std::path::PathBuf;
use rocm_kernel_macros::{amdgpu_kernel_attr, amdgpu_kernel_finalize, amdgpu_kernel_init};
use rocm_rs::hip::*;
use rocm_rs::{
hip::{kernel::AsKernelArg, *},
kernel_args,
rocm_kernel_macros::{amdgpu_global, amdgpu_kernel_finalize, amdgpu_kernel_init},
};
const LEN: usize = 1024;
@@ -59,7 +62,7 @@ const LEN: usize = 1024;
amdgpu_kernel_init!();
// marking code that will be coppied to gpu kernel
#[amdgpu_kernel_attr]
#[amdgpu_global]
fn kernel(input: *const u32, output: *mut u32) {
// retriving data from buffere by workitem
let num = read_by_workitem_id_x(input);
@@ -100,7 +103,7 @@ fn main() -> Result<()> {
input.copy_from_host(&in_host)?;
// providing arguments for kernel
let kernel_args = [input.as_kernel_arg(), output.as_kernel_arg()];
let kernel_args = kernel_args!(input, output);
// setting up launch args
let grid_dim = Dim3 { x: 2, y: 1, z: 1 };
@@ -119,6 +122,64 @@ fn main() -> Result<()> {
Ok(())
}
```
For async operations with streams:
```rust
use std::path::PathBuf;
use rocm_rs::{hip::{kernel::AsKernelArg, *}, rocm_kernel_macros::*};
const LEN: usize = 1024;
amdgpu_kernel_init!();
#[amdgpu_global]
fn kernel(input: *const u32, output: *mut u32) {
let num = read_by_workitem_id_x(input);
write_by_workitem_id_x(output, num * 3);
}
const AMDGPU_KERNEL_BINARY_PATH: &str = amdgpu_kernel_finalize!();
fn main() -> Result<()> {
let device = Device::current()?;
// Create a stream for async operations
let stream = device.get_stream()?;
stream.add_callback(|| println!("callback"))?;
let kernel_path = PathBuf::from(AMDGPU_KERNEL_BINARY_PATH);
let module = Module::load(kernel_path)?;
let function = module.get_function("kernel")?;
let in_host: Vec<u32> = (0..LEN).map(|i| i as u32).collect();
let out_host: Vec<u32> = vec![0; LEN];
let input = DeviceMemory::<u32>::new(LEN)?;
let output = DeviceMemory::<u32>::new(LEN)?;
// Copy data from host to device asynchronously
input.copy_from_host_async(in_host, &stream)?;
let kernel_args = [input.as_kernel_arg(), output.as_kernel_arg()];
let grid_dim = Dim3 { x: 2, y: 1, z: 1 };
let block_dim = Dim3 { x: (LEN / 2) as u32, y: 1, z: 1 };
function.launch(grid_dim, block_dim, 0, Some(&stream), &mut kernel_args.clone())?;
// Retrieve computed data asynchronously
let pending = output.copy_to_host_async(out_host, &stream)?;
// Synchronize memory (awaiting for copy to finish)
let out_host = stream.synchronize_memory(pending)?;
println!("Output: {:?}", &out_host[..256]);
Ok(())
}
```
```
@@ -179,13 +240,28 @@ cargo build
- rocm_smi - enables bindings and wrappers for rocm_smi_lib
## Examples
- hip
- vector_add - example containing kernel written in cpp launched with rocm-rs
- rust_kernel - example containing kernel written in in rust using macros
- rust_kernel_async - example containing kernel written in in rust, using stream to manage memory asynchronously
- saxpy - X = aX+Y
- rand
- normal - generating random numbers with normal distribution
The project includes a workspace with examples for each sub-library. Run examples with:
```bash
cargo run --package <example_name> --example <example_name>
```
### HIP
- `vector_add_example` - Vector addition with kernel written in HIP
- `rust_kernel` - Kernel written in Rust using macros
- `rust_kernel_async` - Async kernel execution with streams
- `saxpy` - SAXPY operation (Single-precision A*X + Y)
- `sort` - GPU sorting example
### MIOpen
- `miopen_basic` - Basic MIOpen usage example
- `multi_tensor` - Multi-tensor operations
### rocBLAS
- `rocblas_basic` - Basic rocBLAS usage example
### rocRAND
- `normal` - Random number generation with normal distribution
## Contributing
src/hip/bindings.rs
+31 -3
View File
@@ -138,12 +138,12 @@ where
}
}
pub const HIP_VERSION_MAJOR: u32 = 7;
pub const HIP_VERSION_MINOR: u32 = 1;
pub const HIP_VERSION_PATCH: u32 = 52802;
pub const HIP_VERSION_MINOR: u32 = 2;
pub const HIP_VERSION_PATCH: u32 = 26064;
pub const HIP_VERSION_GITHASH: &[u8; 1] = b"\0";
pub const HIP_VERSION_BUILD_ID: u32 = 0;
pub const HIP_VERSION_BUILD_NAME: &[u8; 1] = b"\0";
pub const HIP_VERSION: u32 = 70152802;
pub const HIP_VERSION: u32 = 70226064;
pub const HIP_TRSA_OVERRIDE_FORMAT: u32 = 1;
pub const HIP_TRSF_READ_AS_INTEGER: u32 = 1;
pub const HIP_TRSF_NORMALIZED_COORDINATES: u32 = 2;
@@ -1290,6 +1290,7 @@ pub const hipDeviceAttribute_t_hipDeviceAttributeVirtualMemoryManagementSupporte
pub const hipDeviceAttribute_t_hipDeviceAttributeHostRegisterSupported: hipDeviceAttribute_t = 90;
pub const hipDeviceAttribute_t_hipDeviceAttributeMemoryPoolSupportedHandleTypes:
hipDeviceAttribute_t = 91;
pub const hipDeviceAttribute_t_hipDeviceAttributeHostNumaId: hipDeviceAttribute_t = 92;
pub const hipDeviceAttribute_t_hipDeviceAttributeCudaCompatibleEnd: hipDeviceAttribute_t = 9999;
pub const hipDeviceAttribute_t_hipDeviceAttributeAmdSpecificBegin: hipDeviceAttribute_t = 10000;
pub const hipDeviceAttribute_t_hipDeviceAttributeClockInstructionRate: hipDeviceAttribute_t = 10000;
@@ -3152,6 +3153,9 @@ unsafe extern "C" {
value_out: *mut hipLaunchAttributeValue,
) -> hipError_t;
}
unsafe extern "C" {
pub fn hipStreamCopyAttributes(dst: hipStream_t, src: hipStream_t) -> hipError_t;
}
unsafe extern "C" {
pub fn hipStreamWaitValue32(
stream: hipStream_t,
@@ -4306,6 +4310,22 @@ unsafe extern "C" {
library: hipLibrary_t,
) -> hipError_t;
}
unsafe extern "C" {
pub fn hipLibraryEnumerateKernels(
kernels: *mut hipKernel_t,
numKernels: ::std::os::raw::c_uint,
library: hipLibrary_t,
) -> hipError_t;
}
unsafe extern "C" {
pub fn hipKernelGetLibrary(library: *mut hipLibrary_t, kernel: hipKernel_t) -> hipError_t;
}
unsafe extern "C" {
pub fn hipKernelGetName(
name: *mut *const ::std::os::raw::c_char,
kernel: hipKernel_t,
) -> hipError_t;
}
unsafe extern "C" {
pub fn hipFuncGetAttributes(
attr: *mut hipFuncAttributes,
@@ -4541,6 +4561,14 @@ unsafe extern "C" {
blockSizeLimit: ::std::os::raw::c_int,
) -> hipError_t;
}
unsafe extern "C" {
pub fn hipOccupancyAvailableDynamicSMemPerBlock(
dynamicSmemSize: *mut usize,
f: *const ::std::os::raw::c_void,
numBlocks: ::std::os::raw::c_int,
blockSize: ::std::os::raw::c_int,
) -> hipError_t;
}
unsafe extern "C" {
pub fn hipProfilerStart() -> hipError_t;
}
src/hip/examples/rust_kernel/Cargo.lock
Generated
+2 -27
View File
@@ -144,12 +144,6 @@ dependencies = [
"minimal-lexical",
]
[[package]]
name = "once_cell"
version = "1.21.3"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "42f5e15c9953c5e4ccceeb2e7382a716482c34515315f7b03532b8b4e8393d2d"
[[package]]
name = "paste"
version = "1.0.15"
@@ -215,34 +209,16 @@ checksum = "2b15c43186be67a4fd63bee50d0303afffcef381492ebe2c5d87f324e1b8815c"
[[package]]
name = "rocm-rs"
version = "0.4.2"
version = "0.5.0"
dependencies = [
"bindgen",
"log",
"once_cell",
"paste",
"rocm_kernel_macros 0.4.2",
]
[[package]]
name = "rocm_kernel_macros"
version = "0.3.0"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "b2c059e17b93ebabc4502167a09e4c56626d88e5fba6b070785ebd7ce3c387e1"
dependencies = [
"fslock",
"proc-macro2",
"quote",
"serde",
"serde_json",
"syn",
"rocm_kernel_macros",
]
[[package]]
name = "rocm_kernel_macros"
version = "0.4.2"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "c180cfaeffccc518db1100beffb3eae2827ae52ad89c56a5417aa0df37f13ca3"
dependencies = [
"fslock",
"proc-macro2",
@@ -258,7 +234,6 @@ name = "rust_kernel"
version = "0.1.0"
dependencies = [
"rocm-rs",
"rocm_kernel_macros 0.3.0",
]
[[package]]
src/hip/examples/rust_kernel/Cargo.toml
-1
View File
@@ -4,5 +4,4 @@ version = "0.1.0"
edition = "2024"
[dependencies]
rocm_kernel_macros = "0.3.0"
rocm-rs = { path = "../../../.." }
src/hip/examples/rust_kernel/src/main.rs
+2 -3
View File
@@ -1,9 +1,8 @@
use std::path::PathBuf;
use rocm_kernel_macros::{amdgpu_kernel_attr, amdgpu_kernel_finalize, amdgpu_kernel_init};
use rocm_rs::{
hip::{kernel::AsKernelArg, *},
kernel_args,
kernel_args, rocm_kernel_macros::{amdgpu_global, amdgpu_kernel_finalize, amdgpu_kernel_init},
};
const LEN: usize = 1024;
@@ -12,7 +11,7 @@ const LEN: usize = 1024;
amdgpu_kernel_init!();
// marking code that will be coppied to gpu kernel
#[amdgpu_kernel_attr]
#[amdgpu_global]
fn kernel(input: *const u32, output: *mut u32) {
// retriving data from buffere by workitem
let num = read_by_workitem_id_x(input);
src/hip/examples/vector_add/main.rs
+3 -2
View File
@@ -1,4 +1,5 @@
use rocm_rs::error::Result;
use rocm_rs::hip::kernel::AsKernelArg;
use rocm_rs::hip::{calculate_grid_1d, Device, DeviceMemory, Dim3, Module, Stream, Timer};
use std::env;
use std::path::PathBuf;
@@ -39,7 +40,7 @@ fn main() -> Result<()> {
let function = module.get_function("vector_add")?;
// Create a stream for async operations
let stream = Stream::new()?;
let stream = device.get_stream()?;
// Get the test size from command line or use default
let args: Vec<String> = env::args().collect();
@@ -54,7 +55,7 @@ fn main() -> Result<()> {
// Prepare host data
println!("Preparing host data...");
let a: Vec<f32> = (0..n).map(|i| i as f32).collect();
let b: Vec<f32> = (0..n).map(|i| (2.0 * i as f32)).collect();
let b: Vec<f32> = (0..n).map(|i| 2.0 * i as f32).collect();
let c = vec![0.0f32; n];
// Allocate device memory
src/hip/memory_ext/sorting.rs
+18 -15
View File
@@ -3,9 +3,12 @@ use rocm_kernel_macros::{
amdgpu_device, amdgpu_global, amdgpu_kernel_finalize, amdgpu_kernel_init,
};
amdgpu_kernel_init!(path: __build_in_kernels_sorting);
#[amdgpu_device(__build_in_kernels_sorting)]
amdgpu_kernel_init!(
path = __build_in_kernels_sorting,
dir = __build_in_kernels_sorting,
binary_name = sorting
);
#[amdgpu_device(path = __build_in_kernels_sorting, dir = __build_in_kernels_sorting)]
use core::{cmp::PartialOrd, ptr::swap};
use crate::{
@@ -13,7 +16,7 @@ use crate::{
kernel_args,
};
#[amdgpu_device(__build_in_kernels_sorting)]
#[amdgpu_device(path = __build_in_kernels_sorting, dir = __build_in_kernels_sorting)]
fn sort_odd_inner<T: Clone + Copy + PartialOrd>(arr: *mut T, ascending: bool) {
let id_x = workgroup_id_x() as usize;
@@ -30,7 +33,7 @@ fn sort_odd_inner<T: Clone + Copy + PartialOrd>(arr: *mut T, ascending: bool) {
}
}
#[amdgpu_device(__build_in_kernels_sorting)]
#[amdgpu_device(path = __build_in_kernels_sorting, dir = __build_in_kernels_sorting)]
fn sort_even_inner<T: Clone + Copy + PartialOrd>(arr: *mut T, ascending: bool) {
let id_x = workgroup_id_x() as usize;
@@ -47,7 +50,7 @@ fn sort_even_inner<T: Clone + Copy + PartialOrd>(arr: *mut T, ascending: bool) {
}
}
#[amdgpu_device(__build_in_kernels_sorting)]
#[amdgpu_device(path = __build_in_kernels_sorting, dir = __build_in_kernels_sorting)]
fn check_sorted_inner<T: Clone + Copy + PartialOrd>(arr: *mut T, target: *mut bool, size: usize) {
let id_x = workgroup_id_x() as usize;
@@ -68,17 +71,17 @@ fn check_sorted_inner<T: Clone + Copy + PartialOrd>(arr: *mut T, target: *mut bo
macro_rules! sort_fns {
($t:ty) => {
paste::paste! {
#[amdgpu_global(__build_in_kernels_sorting)]
#[amdgpu_global(path = __build_in_kernels_sorting, dir = __build_in_kernels_sorting)]
fn [<sort_odd_$t>](arr: *mut $t, ascending: bool) {
sort_odd_inner::<$t>(arr, ascending)
}
#[amdgpu_global(__build_in_kernels_sorting)]
#[amdgpu_global(path = __build_in_kernels_sorting, dir = __build_in_kernels_sorting)]
fn [<sort_even_$t>](arr: *mut $t, ascending: bool) {
sort_even_inner::<$t>(arr, ascending)
}
#[amdgpu_global(__build_in_kernels_sorting)]
#[amdgpu_global(path = __build_in_kernels_sorting, dir = __build_in_kernels_sorting)]
fn [<check_sorted_$t>](arr: *mut $t, target: *mut bool, size: usize) {
check_sorted_inner::<$t>(arr, target, size)
}
@@ -99,8 +102,11 @@ macro_rules! impl_gpu_sort_allowed {
impl_gpu_sort_allowed!(i8, i16, i32, i64, u8, u16, u32, u64, f32, f64);
pub(crate) const SORTING_KERNEL: &[u8] =
include_bytes!(amdgpu_kernel_finalize!(__build_in_kernels_sorting));
pub(crate) const SORTING_KERNEL: &[u8] = include_bytes!(amdgpu_kernel_finalize!(
path = __build_in_kernels_sorting,
dir = __build_in_kernels_sorting,
binary_name = sorting
));
pub(crate) fn sort<T>(mem: &mut DeviceMemory<T>, stream: &Stream, ascending: bool) -> Result<()> {
let module = Module::load_data(SORTING_KERNEL)?;
@@ -161,10 +167,7 @@ pub(crate) fn check_sorted<T>(mem: &DeviceMemory<T>, stream: Option<&Stream>) ->
mod test {
use crate::{
error::Result,
hip::{
Device, DeviceMemory,
memory_ext::sorting::check_sorted,
},
hip::{Device, DeviceMemory, memory_ext::sorting::check_sorted},
};
#[test]
src/miopen/bindings.rs
+2 -2
View File
@@ -41,8 +41,8 @@ pub const MIOPEN_NDEBUG: u32 = 0;
pub const MIOPEN_ALLOC_BUFFERS: u32 = 0;
pub const MIOPEN_ROCBLAS_VERSION_MAJOR: u32 = 5;
pub const MIOPEN_ROCBLAS_VERSION_MINOR: u32 = 1;
pub const MIOPEN_ROCBLAS_VERSION_PATCH: u32 = 1;
pub const MIOPEN_ROCBLAS_VERSION_FLAT: u32 = 5001001;
pub const MIOPEN_ROCBLAS_VERSION_PATCH: u32 = 0;
pub const MIOPEN_ROCBLAS_VERSION_FLAT: u32 = 5001000;
pub const MIOPEN_GOLDEN_DB_VERSION: u32 = 21;
pub const MIOPEN_API_VERSION_REDUCE_TENSOR: u32 = 1;
#[repr(C)]
src/miopen/examples/basic/Cargo.toml
+1 -1
View File
@@ -1,5 +1,5 @@
[package]
name = "basic"
name = "miopen_basic"
version = "0.1.0"
edition = "2024"
src/rocblas/examples/basic/Cargo.toml
+1 -1
View File
@@ -1,5 +1,5 @@
[package]
name = "basic"
name = "rocblas_basic"
version = "0.1.0"
edition = "2024"