91d4e40e02
Add a rustc intrinsic `amdgpu_dispatch_ptr` to access the kernel dispatch packet on amdgpu. The HSA kernel dispatch packet contains important information like the launch size and workgroup size. The Rust intrinsic lowers to the `llvm.amdgcn.dispatch.ptr` LLVM intrinsic, which returns a `ptr addrspace(4)`, plus an addrspacecast to `addrspace(0)`, so it can be returned as a Rust reference. The returned pointer/reference is valid for the whole program lifetime, and is therefore `'static`. The return type of the intrinsic (`*const ()`) does not mention the struct so that rustc does not need to know the exact struct type. An alternative would be to define the struct as lang item or add a generic argument to the function. Short version: ```rust #[cfg(target_arch = "amdgpu")] pub fn amdgpu_dispatch_ptr() -> *const (); ```
27 lines
800 B
Rust
27 lines
800 B
Rust
// Tests the amdgpu_dispatch_ptr intrinsic.
|
|
|
|
//@ compile-flags: --crate-type=rlib --target amdgcn-amd-amdhsa -Ctarget-cpu=gfx900
|
|
//@ needs-llvm-components: amdgpu
|
|
//@ add-minicore
|
|
#![feature(intrinsics, no_core, rustc_attrs)]
|
|
#![no_core]
|
|
|
|
extern crate minicore;
|
|
|
|
pub struct DispatchPacket {
|
|
pub header: u16,
|
|
pub setup: u16,
|
|
pub workgroup_size_x: u16, // and more
|
|
}
|
|
|
|
#[rustc_intrinsic]
|
|
#[rustc_nounwind]
|
|
fn amdgpu_dispatch_ptr() -> *const ();
|
|
|
|
// CHECK-LABEL: @get_dispatch_data
|
|
// CHECK: %[[ORIG_PTR:[^ ]+]] = {{(tail )?}}call ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
|
|
// CHECK-NEXT: %[[PTR:[^ ]+]] = addrspacecast ptr addrspace(4) %[[ORIG_PTR]] to ptr
|
|
#[unsafe(no_mangle)]
|
|
pub fn get_dispatch_data() -> &'static DispatchPacket {
|
|
unsafe { &*(amdgpu_dispatch_ptr() as *const _) }
|
|
}
|