diff --git a/.gitmodules b/.gitmodules new file mode 100644 index 0000000..d76dec3 --- /dev/null +++ b/.gitmodules @@ -0,0 +1,4 @@ +[submodule "crates/rustc_offload_frontend"] + path = crates/rustc_offload_frontend + url = https://github.com/sa4dus/rustc_offload_frontend + branch = main diff --git a/Cargo.lock b/Cargo.lock index 0f5194d..3546f3d 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -4,13 +4,21 @@ version = 4 [[package]] name = "libc" -version = "0.2.177" +version = "0.2.186" source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "2874a2af47a2325c2001a6e6fad9b16a53b802102b528163885171cf92b15976" +checksum = "68ab91017fe16c622486840e4c83c9a37afeff978bd239b5293d61ece587de66" [[package]] name = "rust_perf" version = "0.1.0" +dependencies = [ + "libc", + "rustc_offload_frontend", +] + +[[package]] +name = "rustc_offload_frontend" +version = "0.1.0" dependencies = [ "libc", ] diff --git a/Cargo.toml b/Cargo.toml index c1032a5..e1c3398 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -1,3 +1,5 @@ +cargo-features = ["panic-immediate-abort"] + [package] name = "rust_perf" version = "0.1.0" @@ -5,9 +7,22 @@ edition = "2024" [lib] crate-type = ["cdylib"] +#crate-type = ["rlib", "cdylib"] +path = "src/lib.rs" [dependencies] libc = { version = "0.2.175", default-features = false } +rustc_offload_frontend = { path = "crates/rustc_offload_frontend" } + +[profile.release] +lto = "fat" +panic = "immediate-abort" +codegen-units = 1 + +[profile.dev] +lto = "fat" +panic = "immediate-abort" + [features] default = ["all", "f64"] @@ -40,10 +55,3 @@ vol3d = [] f64 = [] a64 = [] -[profile.release] -lto = "fat" -panic = "abort" - -[profile.dev] -lto = "fat" -panic = "abort" diff --git a/crates/rustc_offload_frontend/.gitignore b/crates/rustc_offload_frontend/.gitignore new file mode 100644 index 0000000..ea8c4bf --- /dev/null +++ b/crates/rustc_offload_frontend/.gitignore @@ -0,0 +1 @@ +/target diff --git a/crates/rustc_offload_frontend/Cargo.lock b/crates/rustc_offload_frontend/Cargo.lock new file mode 100644 index 0000000..c4532ca --- /dev/null +++ b/crates/rustc_offload_frontend/Cargo.lock @@ -0,0 +1,16 @@ +# This file is automatically @generated by Cargo. +# It is not intended for manual editing. +version = 4 + +[[package]] +name = "libc" +version = "0.2.186" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "68ab91017fe16c622486840e4c83c9a37afeff978bd239b5293d61ece587de66" + +[[package]] +name = "rustc_offload_frontend" +version = "0.1.0" +dependencies = [ + "libc", +] diff --git a/crates/rustc_offload_frontend/Cargo.toml b/crates/rustc_offload_frontend/Cargo.toml new file mode 100644 index 0000000..d672440 --- /dev/null +++ b/crates/rustc_offload_frontend/Cargo.toml @@ -0,0 +1,30 @@ +[package] +name = "rustc_offload_frontend" +version = "0.1.0" +edition = "2024" +license = "MIT" +description = "rustc offload feature frontend draft" +repository = "https://github.com/sa4dus/rustc_offload_frontend" +readme = "README.md" + +[lib] +crate-type = ["rlib"] +#crate-type = ["rlib", "cdylib"] +path = "src/lib.rs" + +[[bin]] +name = "main" +path = "src/main.rs" + +[dependencies] +libc = { version = "0.2.175", default-features = false } + +[features] + +[profile.release] +lto = "fat" +panic = "abort" + +[profile.dev] +lto = "fat" +panic = "abort" diff --git a/crates/rustc_offload_frontend/README.md b/crates/rustc_offload_frontend/README.md new file mode 100644 index 0000000..d518681 --- /dev/null +++ b/crates/rustc_offload_frontend/README.md @@ -0,0 +1 @@ +# rustc_offload_frontend diff --git a/crates/rustc_offload_frontend/src/gpu.rs b/crates/rustc_offload_frontend/src/gpu.rs new file mode 100644 index 0000000..46c77ea --- /dev/null +++ b/crates/rustc_offload_frontend/src/gpu.rs @@ -0,0 +1,191 @@ +#[derive(Clone, Copy)] +pub struct Dim3 { + pub x: usize, + pub y: usize, + pub z: usize, +} + +#[cfg(target_arch = "amdgpu")] +use core::intrinsics::gpu::amdgpu_dispatch_ptr; + +/// Get the packet for this dispatch. +/// +/// Get a reference to the packet that was used to dispatch this kernel. +/// The dispatch packet contains information like the workgroup size and dispatch size. +/// +/// # Example +/// +/// ```rust +/// # #![no_std] +/// # extern crate alloc; +/// # fn main() { +/// use amdgpu_device_libs::prelude::*; +/// +/// let dispatch = dispatch_ptr(); +/// println!("Workgroup size {}x{}x{}", dispatch.workgroup_size_x, dispatch.workgroup_size_y, dispatch.workgroup_size_z); +/// # } +/// ``` +#[cfg(target_arch = "amdgpu")] +#[inline] +pub fn dispatch_ptr() -> &'static HsaKernelDispatchPacket { + unsafe { + &*core::mem::transmute::<*const (), *const HsaKernelDispatchPacket>( + core::intrinsics::gpu::amdgpu_dispatch_ptr(), + ) + } +} +/// HSA packet to dispatch a kernel. +/// +/// A pointer to the packet that was used to dispatch the currently running kernel can be obtained with [`dispatch_ptr`]. +#[derive(Clone, Eq, PartialEq, Ord, PartialOrd, Hash)] +#[repr(C)] +pub struct HsaKernelDispatchPacket { + /// Packet header. Used to configure multiple packet parameters such as the + /// packet type. The parameters are described by hsa_packet_header_t. + pub header: u16, + /// Dispatch setup parameters. Used to configure kernel dispatch parameters + /// such as the number of dimensions in the grid. The parameters are described + /// by hsa_kernel_dispatch_packet_setup_t. + pub setup: u16, + /// X dimension of work-group, in work-items. Must be greater than 0. + pub workgroup_size_x: u16, + /// Y dimension of work-group, in work-items. Must be greater than + /// 0. If the grid has 1 dimension, the only valid value is 1. + pub workgroup_size_y: u16, + /// Z dimension of work-group, in work-items. Must be greater than + /// 0. If the grid has 1 or 2 dimensions, the only valid value is 1. + pub workgroup_size_z: u16, + /// Reserved. Must be 0. + pub reserved0: u16, + /// X dimension of grid, in work-items. Must be greater than 0. Must + /// not be smaller than @a workgroup_size_x. + pub grid_size_x: u32, + /// Y dimension of grid, in work-items. Must be greater than 0. If the grid has + /// 1 dimension, the only valid value is 1. Must not be smaller than @a + /// workgroup_size_y. + pub grid_size_y: u32, + /// Z dimension of grid, in work-items. Must be greater than 0. If the grid has + /// 1 or 2 dimensions, the only valid value is 1. Must not be smaller than @a + /// workgroup_size_z. + pub grid_size_z: u32, + /// Size in bytes of private memory allocation request (per work-item). + pub private_segment_size: u32, + /// Size in bytes of group memory allocation request (per work-group). Must not + /// be less than the sum of the group memory used by the kernel (and the + /// functions it calls directly or indirectly) and the dynamically allocated + /// group segment variables. + pub group_segment_size: u32, + /// Opaque handle to a code object that includes an implementation-defined + /// executable code for the kernel. + pub kernel_object: u64, + /// Pointer to the kernel arguments. + pub kernarg_address: *mut core::ffi::c_void, + /// Reserved. Must be 0. + pub reserved2: u64, + /// Signal used to indicate completion of the job. The application can use the + /// special signal handle 0 to indicate that no signal is used. + pub completion_signal: u64, +} +// Handle to an HSA signal. +//#[cfg(feature = "device_libs")] +//#[derive(Copy, Clone, Eq, PartialEq, Ord, PartialOrd, Hash)] +//#[repr(C)] +//pub struct HsaSignal { +// /// The internal representation of an HSA signal. +// pub handle: u64, +//} + +pub(crate) fn global_thread_dim() -> Dim3 { + #[cfg(target_arch = "nvptx64")] + unsafe { + use core::arch::nvptx::*; + Dim3 { + x: (_block_idx_x() * _block_dim_x() + _thread_idx_x()) as usize, + y: (_block_idx_y() * _block_dim_y() + _thread_idx_y()) as usize, + z: (_block_idx_z() * _block_dim_z() + _thread_idx_z()) as usize, + } + } + #[cfg(target_arch = "amdgpu")] + unsafe { + use core::arch::amdgpu::*; + let dispatch = dispatch_ptr(); + + let x = (workgroup_id_x() * (*dispatch).workgroup_size_x as u32 + workitem_id_x()) as usize; + let y = (workgroup_id_y() * (*dispatch).workgroup_size_y as u32 + workitem_id_y()) as usize; + let z = (workgroup_id_z() * (*dispatch).workgroup_size_z as u32 + workitem_id_z()) as usize; + Dim3 { x, y, z } + } + #[cfg(target_os = "linux")] + Dim3 { x: 0, y: 0, z: 0 } +} + +pub(crate) fn block_idx() -> Dim3 { + #[cfg(target_arch = "nvptx64")] + unsafe { + use core::arch::nvptx::*; + Dim3 { + x: _block_idx_x() as usize, + y: _block_idx_y() as usize, + z: _block_idx_z() as usize, + } + } + #[cfg(target_arch = "amdgpu")] + unsafe { + use core::arch::amdgpu::*; + let dispatch = dispatch_ptr(); + + let x = (workgroup_id_x()) as usize; + let y = (workgroup_id_y()) as usize; + let z = (workgroup_id_z()) as usize; + Dim3 { x, y, z } + } + #[cfg(target_os = "linux")] + Dim3 { x: 0, y: 0, z: 0 } +} + +pub(crate) fn block_dim() -> Dim3 { + #[cfg(target_arch = "nvptx64")] + unsafe { + use core::arch::nvptx::*; + Dim3 { + x: _block_dim_x() as usize, + y: _block_dim_y() as usize, + z: _block_dim_z() as usize, + } + } + #[cfg(target_arch = "amdgpu")] + unsafe { + use core::arch::amdgpu::*; + let dispatch = dispatch_ptr(); + + let x = dispatch.workgroup_size_x as usize; + let y = dispatch.workgroup_size_y as usize; + let z = dispatch.workgroup_size_z as usize; + Dim3 { x, y, z } + } + #[cfg(target_os = "linux")] + Dim3 { x: 0, y: 0, z: 0 } +} + +pub(crate) fn thread_idx() -> Dim3 { + #[cfg(target_arch = "nvptx64")] + unsafe { + use core::arch::nvptx::*; + Dim3 { + x: _thread_idx_x() as usize, + y: _thread_idx_y() as usize, + z: _thread_idx_z() as usize, + } + } + #[cfg(target_arch = "amdgpu")] + unsafe { + use core::arch::amdgpu::*; + + let x = (workitem_id_x()) as usize; + let y = (workitem_id_y()) as usize; + let z = (workitem_id_z()) as usize; + Dim3 { x, y, z } + } + #[cfg(target_os = "linux")] + Dim3 { x: 0, y: 0, z: 0 } +} diff --git a/crates/rustc_offload_frontend/src/lib.rs b/crates/rustc_offload_frontend/src/lib.rs new file mode 100644 index 0000000..9ed46b7 --- /dev/null +++ b/crates/rustc_offload_frontend/src/lib.rs @@ -0,0 +1,70 @@ +#![allow(internal_features)] +#![allow(linker_messages)] +#![allow(improper_ctypes)] +#![allow(improper_gpu_kernel_arg)] +#![allow(improper_ctypes_definitions)] +#![feature(gpu_offload, core_intrinsics, gpu_intrinsics, offload)] +#![cfg_attr(target_arch = "nvptx64", feature(stdarch_nvptx))] +#![cfg_attr(target_arch = "amdgpu", feature(stdarch_amdgpu))] +#![cfg_attr(any(target_arch = "nvptx64", target_arch = "amdgpu"), no_std)] + +pub use core::offload::offload_kernel; + +pub mod gpu; +pub mod partition; + +#[macro_export] +macro_rules! offload { + ( $($field:ident = $val:expr),* $(,)? ) => { + $crate::offload!(@munch + [ $($field = $val),* ]; + kernel = NONE; + grid_dim = ([1, 1, 1]); + block_dim = ([1, 1, 1]); + dyn_cache = (0); + args = NONE + ); + }; + + (@munch [kernel = $val:expr $(, $rest_f:ident = $rest_v:expr)*]; kernel = $k:tt; grid_dim = $g:tt; block_dim = $b:tt; dyn_cache = $d:tt; args = $a:tt) => { + $crate::offload!(@munch [$($rest_f = $rest_v),*]; kernel = (SOME $val); grid_dim = $g; block_dim = $b; dyn_cache = $d; args = $a); + }; + (@munch [grid_dim = $val:expr $(, $rest_f:ident = $rest_v:expr)*]; kernel = $k:tt; grid_dim = $g:tt; block_dim = $b:tt; dyn_cache = $d:tt; args = $a:tt) => { + $crate::offload!(@munch [$($rest_f = $rest_v),*]; kernel = $k; grid_dim = ($val); block_dim = $b; dyn_cache = $d; args = $a); + }; + (@munch [block_dim = $val:expr $(, $rest_f:ident = $rest_v:expr)*]; kernel = $k:tt; grid_dim = $g:tt; block_dim = $b:tt; dyn_cache = $d:tt; args = $a:tt) => { + $crate::offload!(@munch [$($rest_f = $rest_v),*]; kernel = $k; grid_dim = $g; block_dim = ($val); dyn_cache = $d; args = $a); + }; + (@munch [dyn_cache = $val:expr $(, $rest_f:ident = $rest_v:expr)*]; kernel = $k:tt; grid_dim = $g:tt; block_dim = $b:tt; dyn_cache = $d:tt; args = $a:tt) => { + $crate::offload!(@munch [$($rest_f = $rest_v),*]; kernel = $k; grid_dim = $g; block_dim = $b; dyn_cache = ($val); args = $a); + }; + (@munch [args = $val:expr $(, $rest_f:ident = $rest_v:expr)*]; kernel = $k:tt; grid_dim = $g:tt; block_dim = $b:tt; dyn_cache = $d:tt; args = $a:tt) => { + $crate::offload!(@munch [$($rest_f = $rest_v),*]; kernel = $k; grid_dim = $g; block_dim = $b; dyn_cache = $d; args = (SOME $val)); + }; + + (@munch [$invalid:ident = $val:expr $(, $rest_f:ident = $rest_v:expr)*]; kernel = $k:tt; grid_dim = $g:tt; block_dim = $b:tt; dyn_cache = $d:tt; args = $a:tt) => { + compile_error!(concat!("unknown field ", stringify!($invalid))); + }; + + (@munch []; kernel = NONE; grid_dim = $g:tt; block_dim = $b:tt; dyn_cache = $d:tt; args = $a:tt) => { + compile_error!("missing `kernel`"); + }; + (@munch []; kernel = $k:tt; grid_dim = $g:tt; block_dim = $b:tt; dyn_cache = $d:tt; args = NONE) => { + compile_error!("missing `args`"); + }; + (@munch []; kernel = (SOME $kernel:expr); grid_dim = ($grid_dim:expr); block_dim = ($block_dim:expr); dyn_cache = ($dyn_cache:expr); args = (SOME $args:expr)) => { + core::intrinsics::offload::<_, _, ()>( + $kernel, + $grid_dim, + $block_dim, + $dyn_cache, + $args, + ) + }; +} + +#[cfg(target_arch = "nvptx64")] +#[panic_handler] +fn panic(_: &core::panic::PanicInfo) -> ! { + loop {} +} diff --git a/crates/rustc_offload_frontend/src/partition.rs b/crates/rustc_offload_frontend/src/partition.rs new file mode 100644 index 0000000..bdeaa40 --- /dev/null +++ b/crates/rustc_offload_frontend/src/partition.rs @@ -0,0 +1,323 @@ +use crate::gpu::{block_dim, block_idx, global_thread_dim, thread_idx}; +use core::convert::From; +use core::offload::offload::PreloadMut; +use core::prelude::v1::*; + +pub unsafe trait PartitioningStrategy { + type View<'a, T: 'a>; + type ViewMut<'a, T: 'a>; + + fn index() -> usize; + unsafe fn get<'a, T>(ptr: *const T, len: usize) -> Option>; + unsafe fn get_mut<'a, T>(ptr: *mut T, len: usize) -> Option>; +} + +#[derive(Copy, Clone)] +pub struct Region<'a, T, S: PartitioningStrategy> { + ptr: *mut T, + len: usize, + _marker: core::marker::PhantomData<(&'a mut [T], S)>, +} + +impl<'a, T, const N: usize, S> From<&PreloadMut<'a, [T; N]>> for Region<'a, T, S> +where + S: PartitioningStrategy, +{ + fn from(p: &PreloadMut<'a, [T; N]>) -> Self { + Self { + ptr: p.cpu_ptr as *mut T, + len: N, + _marker: core::marker::PhantomData, + } + } +} + +pub struct RawRegion<'a, T> { + pub ptr: *mut T, + pub len: usize, + _marker: core::marker::PhantomData<&'a mut [T]>, +} + +impl<'a, T> From<&'a mut [T]> for RawRegion<'a, T> { + fn from(data: &'a mut [T]) -> Self { + Self { + ptr: data.as_mut_ptr(), + len: data.len(), + _marker: core::marker::PhantomData, + } + } +} + +impl<'a, T, const N: usize> From<&'a mut [T; N]> for RawRegion<'a, T> { + fn from(data: &'a mut [T; N]) -> Self { + Self { + ptr: data.as_mut_ptr(), + len: N, + _marker: core::marker::PhantomData, + } + } +} + +impl<'a, T, S: PartitioningStrategy> Region<'a, T, S> { + pub fn new(data: D) -> Self + where + D: Into>, + { + let raw = data.into(); + Self { + ptr: raw.ptr, + len: raw.len, + _marker: core::marker::PhantomData, + } + } + + pub fn get(&self) -> Option> { + unsafe { S::get(self.ptr as *const T, self.len) } + } + + pub fn get_mut(&mut self) -> Option> { + unsafe { S::get_mut(self.ptr, self.len) } + } +} + +// linear1d +#[derive(Copy, Clone)] +pub struct Linear1D; +unsafe impl PartitioningStrategy for Linear1D { + type View<'a, T: 'a> = &'a T; + type ViewMut<'a, T: 'a> = &'a mut T; + + fn index() -> usize { + global_thread_dim().x + } + unsafe fn get<'a, T>(ptr: *const T, len: usize) -> Option> { + let idx = Self::index(); + if idx < len { + Some(unsafe { &*ptr.add(idx) }) + } else { + None + } + } + unsafe fn get_mut<'a, T>(ptr: *mut T, len: usize) -> Option> { + let idx = Self::index(); + if idx < len { + Some(unsafe { &mut *ptr.add(idx) }) + } else { + None + } + } +} + +// linear2d +#[derive(Copy, Clone)] +pub struct Linear2D; +unsafe impl PartitioningStrategy for Linear2D { + type View<'a, T: 'a> = &'a T; + type ViewMut<'a, T: 'a> = &'a mut T; + + fn index() -> usize { + let tid = global_thread_dim(); + tid.y * W + tid.x + } + unsafe fn get<'a, T>(ptr: *const T, len: usize) -> Option> { + let idx = Self::index(); + if idx < len { + Some(unsafe { &*ptr.add(idx) }) + } else { + None + } + } + unsafe fn get_mut<'a, T>(ptr: *mut T, len: usize) -> Option> { + let idx = Self::index(); + if idx < len { + Some(unsafe { &mut *ptr.add(idx) }) + } else { + None + } + } +} + +// stride1d +#[derive(Copy, Clone)] +pub struct Stride1D; +unsafe impl PartitioningStrategy for Stride1D { + type View<'a, T: 'a> = &'a T; + type ViewMut<'a, T: 'a> = &'a mut T; + + fn index() -> usize { + let bidx = block_idx().x; + let tidx = thread_idx().x; + bidx * STRIDE + tidx + } + unsafe fn get<'a, T>(ptr: *const T, len: usize) -> Option> { + let idx = Self::index(); + if idx < len { + Some(unsafe { &*ptr.add(idx) }) + } else { + None + } + } + unsafe fn get_mut<'a, T>(ptr: *mut T, len: usize) -> Option> { + let idx = Self::index(); + if idx < len { + Some(unsafe { &mut *ptr.add(idx) }) + } else { + None + } + } +} + +// stride2d +#[derive(Copy, Clone)] +pub struct StrideViewMut<'a, T> { + block_ptr: *mut T, + stride: usize, + _marker: core::marker::PhantomData<&'a mut T>, +} +impl<'a, T> StrideViewMut<'a, T> { + pub fn set(&mut self, x: usize, y: usize, val: T) { + unsafe { + *self.block_ptr.add(y * self.stride + x) = val; + } + } +} + +#[derive(Copy, Clone)] +pub struct Stride2D< + const W: usize, + const H: usize, + const SX: usize, + const SY: usize, + const STRIDE: usize, +>; +unsafe impl + PartitioningStrategy for Stride2D +{ + type View<'a, T: 'a> = &'a T; + type ViewMut<'a, T: 'a> = StrideViewMut<'a, T>; + + fn index() -> usize { + let tid = global_thread_dim(); + tid.y * SY * STRIDE + tid.x * SX + } + unsafe fn get<'a, T>(_: *const T, _: usize) -> Option> { + unimplemented!() + } + unsafe fn get_mut<'a, T>(ptr: *mut T, _: usize) -> Option> { + let idx = Self::index(); + Some(StrideViewMut { + block_ptr: unsafe { ptr.add(idx) }, + stride: STRIDE, + _marker: core::marker::PhantomData, + }) + } +} + +// some custom patterns needed for `rust_perf` + +// for vol3d +#[derive(Copy, Clone)] +pub struct OffsetStrideViewMut<'a, T> { + base_ptr: *mut T, + idx: usize, + len: usize, + _marker: core::marker::PhantomData<&'a mut T>, +} + +impl<'a, T> OffsetStrideViewMut<'a, T> { + pub fn set(&mut self, offset: usize, val: T) { + if let Some(final_idx) = self.idx.checked_add(offset) { + if final_idx < self.len { + unsafe { + *self.base_ptr.add(final_idx) = val; + } + } + } + } +} + +#[derive(Copy, Clone)] +pub struct OffsetStride1D; + +unsafe impl PartitioningStrategy for OffsetStride1D { + type View<'a, T: 'a> = &'a T; + type ViewMut<'a, T: 'a> = OffsetStrideViewMut<'a, T>; + + fn index() -> usize { + let bidx = block_idx().x; + let tidx = thread_idx().x; + bidx * STRIDE + tidx + } + + unsafe fn get<'a, T>(_: *const T, _: usize) -> Option> { + unimplemented!("write only") + } + + unsafe fn get_mut<'a, T>(ptr: *mut T, len: usize) -> Option> { + let idx = Self::index(); + + if idx < len { + Some(OffsetStrideViewMut { + base_ptr: ptr, + idx, + len, + _marker: core::marker::PhantomData, + }) + } else { + None + } + } +} + +// for ltimes +#[derive(Copy, Clone)] +pub struct Stride3D< + const BX: usize, + const BY: usize, + const BZ: usize, + const MAX_X: usize, + const MAX_Y: usize, +>; + +unsafe impl + PartitioningStrategy for Stride3D +{ + type View<'a, T: 'a> = &'a T; + type ViewMut<'a, T: 'a> = &'a mut T; + + fn index() -> usize { + let mx = (block_idx().x * BX) + thread_idx().x; + let gy = (block_idx().y * BY) + thread_idx().y; + let zz = (block_idx().z * BZ) + thread_idx().z; + + mx + MAX_X * (gy + MAX_Y * zz) + } + + unsafe fn get<'a, T>(ptr: *const T, len: usize) -> Option> { + let mx = (block_idx().x * BX) + thread_idx().x; + let gy = (block_idx().y * BY) + thread_idx().y; + let zz = (block_idx().z * BZ) + thread_idx().z; + + if mx < MAX_X && gy < MAX_Y { + let idx = mx + MAX_X * (gy + MAX_Y * zz); + if idx < len { + return Some(unsafe { &*ptr.add(idx) }); + } + } + None + } + + unsafe fn get_mut<'a, T>(ptr: *mut T, len: usize) -> Option> { + let mx = (block_idx().x * BX) + thread_idx().x; + let gy = (block_idx().y * BY) + thread_idx().y; + let zz = (block_idx().z * BZ) + thread_idx().z; + + if mx < MAX_X && gy < MAX_Y { + let idx = mx + MAX_X * (gy + MAX_Y * zz); + if idx < len { + return Some(unsafe { &mut *ptr.add(idx) }); + } + } + None + } +} diff --git a/scripts/amd.sh b/scripts/amd.sh new file mode 100755 index 0000000..44cc5c7 --- /dev/null +++ b/scripts/amd.sh @@ -0,0 +1,18 @@ +set -e + +FEATURES_ARGS=("$@") + + +CLANG_LINKER_WRAPPER="/tmp/drehwald1/prog/rust/build/x86_64-unknown-linux-gnu/llvm/bin/clang-linker-wrapper" +#LD_LIBRARY_PATH="/home/mdominguez/workspace/rust/build/x86_64-unknown-linux-gnu/stage1/lib:/home/mdominguez/workspace/rust/build/x86_64-unknown-linux-gnu/offload/lib:$LD_LIBRARY_PATH" + +RUSTFLAGS="-Ctarget-cpu=gfx90a --emit=llvm-bc,llvm-ir -Zoffload=Device -Csave-temps -Zunstable-options" \ +cargo +offload build -Zunstable-options -r -v --target amdgcn-amd-amdhsa -Zbuild-std=core,compiler_builtins,panic_abort -Zbuild-std-features=compiler-builtins-mem "${FEATURES_ARGS[@]}" + +RUSTFLAGS="--emit=llvm-bc,llvm-ir -Csave-temps -Zoffload=Host=/p/lustre1/drehwald1/prog/offload/rust_perf/target/amdgcn-amd-amdhsa/release/deps/device.bin -Zunstable-options" \ +cargo +offload build -r -Zbuild-std=std,panic_abort "${FEATURES_ARGS[@]}" + +DEBUG="" +#"-Xarch_device -mllvm=-print-before-all -Xarch_device -mllvm=-print-module-scope" + +$CLANG_LINKER_WRAPPER "--should-extract=gfx90a" "${DEBUG}" "--device-compiler=amdgcn-amd-amdhsa=-g" "--device-compiler=amdgcn-amd-amdhsa=-save-temps=cwd" "--host-triple=x86_64-unknown-linux-gnu" "--save-temps" "--linker-path=/tmp/drehwald1/prog/rust/build/x86_64-unknown-linux-gnu/lld/bin/ld.lld" "--hash-style=gnu" "--eh-frame-hdr" "-m" "elf_x86_64" "-pie" "-dynamic-linker" "/lib64/ld-linux-x86-64.so.2" "-o" "bare" "/lib/../lib64/Scrt1.o" "/lib/../lib64/crti.o" "/opt/rh/gcc-toolset-12/root/usr/lib/gcc/x86_64-redhat-linux/12/crtbeginS.o" "-L/tmp/drehwald1/prog/rust/build/x86_64-unknown-linux-gnu/llvm/bin/../lib/x86_64-unknown-linux-gnu" "-L/tmp/drehwald1/prog/rust/build/x86_64-unknown-linux-gnu/llvm/lib/clang/21/lib/x86_64-unknown-linux-gnu" "-L/opt/rh/gcc-toolset-12/root/usr/lib/gcc/x86_64-redhat-linux/12" "-L/opt/rh/gcc-toolset-12/root/usr/lib/gcc/x86_64-redhat-linux/12/../../../../lib64" "-L/lib/../lib64" "-L/usr/lib64" "-L/lib" "-L/usr/lib" "./target/release/deps/host.o" "-lstdc++" "-lm" "-L/tmp/drehwald1/prog/rust/build/x86_64-unknown-linux-gnu/offload/lib" "-lomp" "-lomptarget" "-L/tmp/drehwald1/prog/rust/build/x86_64-unknown-linux-gnu/llvm/lib" "-lgcc_s" "-lgcc" "-lpthread" "-lc" "-lgcc_s" "-lgcc" "/opt/rh/gcc-toolset-12/root/usr/lib/gcc/x86_64-redhat-linux/12/crtendS.o" "/lib/../lib64/crtn.o" diff --git a/scripts/compile.sh b/scripts/compile.sh new file mode 100755 index 0000000..ab6a2d5 --- /dev/null +++ b/scripts/compile.sh @@ -0,0 +1,15 @@ +#set -e + +FEATURES_ARGS=("$@") + +CLANG_LINKER_WRAPPER="/tmp/drehwald1/prog/rust/build/x86_64-unknown-linux-gnu/llvm/bin/clang-linker-wrapper" +#LD_LIBRARY_PATH="/home/mdominguez/workspace/rust/build/x86_64-unknown-linux-gnu/stage1/lib:/home/mdominguez/workspace/rust/build/x86_64-unknown-linux-gnu/offload/lib:$LD_LIBRARY_PATH" + +RUSTFLAGS="-Ctarget-cpu=gfx90a --emit=llvm-bc,llvm-ir -Zoffload=Device -Csave-temps -Zunstable-options" \ +cargo +offload build -Zunstable-options -r -v --target amdgcn-amd-amdhsa -Zbuild-std=core $FEATURES_ARGS + +RUSTFLAGS="--emit=llvm-bc,llvm-ir -Csave-temps -Zoffload=Host=/p/lustre1/drehwald1/prog/offload/rust_perf/target/amdgcn-amd-amdhsa/release/deps/device.bin -Zunstable-options" \ +cargo +offload build -r $FEATURES_ARGS + + +$CLANG_LINKER_WRAPPER "--should-extract=gfx90a" "--device-compiler=amdgcn-amd-amdhsa=-g" "--device-compiler=amdgcn-amd-amdhsa=-save-temps=cwd" "--host-triple=x86_64-unknown-linux-gnu" "--save-temps" "--linker-path=/tmp/drehwald1/prog/rust/build/x86_64-unknown-linux-gnu/lld/bin/ld.lld" "--hash-style=gnu" "--eh-frame-hdr" "-m" "elf_x86_64" "-pie" "-dynamic-linker" "/lib64/ld-linux-x86-64.so.2" "-o" "bare" "/lib/../lib64/Scrt1.o" "/lib/../lib64/crti.o" "/opt/rh/gcc-toolset-12/root/usr/lib/gcc/x86_64-redhat-linux/12/crtbeginS.o" "-L/tmp/drehwald1/prog/rust/build/x86_64-unknown-linux-gnu/llvm/bin/../lib/x86_64-unknown-linux-gnu" "-L/tmp/drehwald1/prog/rust/build/x86_64-unknown-linux-gnu/llvm/lib/clang/21/lib/x86_64-unknown-linux-gnu" "-L/opt/rh/gcc-toolset-12/root/usr/lib/gcc/x86_64-redhat-linux/12" "-L/opt/rh/gcc-toolset-12/root/usr/lib/gcc/x86_64-redhat-linux/12/../../../../lib64" "-L/lib/../lib64" "-L/usr/lib64" "-L/lib" "-L/usr/lib" "./target/release/deps/host.o" "-lstdc++" "-lm" "-L/tmp/drehwald1/prog/rust/build/x86_64-unknown-linux-gnu/offload/lib" "-lomp" "-lomptarget" "-L/tmp/drehwald1/prog/rust/build/x86_64-unknown-linux-gnu/llvm/lib" "-lgcc_s" "-lgcc" "-lpthread" "-lc" "-lgcc_s" "-lgcc" "/opt/rh/gcc-toolset-12/root/usr/lib/gcc/x86_64-redhat-linux/12/crtendS.o" "/lib/../lib64/crtn.o" diff --git a/scripts/nv.sh b/scripts/nv.sh new file mode 100755 index 0000000..45f5aef --- /dev/null +++ b/scripts/nv.sh @@ -0,0 +1,44 @@ +#set -e + +FEATURES_ARGS=("$@") + +PTX80_ARGS=( + "--device-compiler=nvptx64-nvidia-cuda=-Xclang" + "--device-compiler=nvptx64-nvidia-cuda=-target-cpu" + "--device-compiler=nvptx64-nvidia-cuda=-Xclang" + "--device-compiler=nvptx64-nvidia-cuda=sm_90a" + "--device-compiler=nvptx64-nvidia-cuda=-Xclang" + "--device-compiler=nvptx64-nvidia-cuda=-target-feature" + "--device-compiler=nvptx64-nvidia-cuda=-Xclang" + "--device-compiler=nvptx64-nvidia-cuda=+ptx80" + + "--device-compiler=nvptx-nvidia-cuda=-Xclang" + "--device-compiler=nvptx-nvidia-cuda=-target-cpu" + "--device-compiler=nvptx-nvidia-cuda=-Xclang" + "--device-compiler=nvptx-nvidia-cuda=sm_90a" + "--device-compiler=nvptx-nvidia-cuda=-Xclang" + "--device-compiler=nvptx-nvidia-cuda=-target-feature" + "--device-compiler=nvptx-nvidia-cuda=-Xclang" + "--device-compiler=nvptx-nvidia-cuda=+ptx80" + "--device-linker=nvptx64-nvidia-cuda=--plugin-opt=-mattr=+ptx80" + "--device-linker=nvptx-nvidia-cuda=--plugin-opt=-mattr=+ptx80" +) + +CLANG_LINKER_WRAPPER="/tmp/drehwald1/prog/rust/build/x86_64-unknown-linux-gnu/llvm/bin/clang-linker-wrapper" +#LD_LIBRARY_PATH="/home/mdominguez/workspace/rust/build/x86_64-unknown-linux-gnu/stage1/lib:/home/mdominguez/workspace/rust/build/x86_64-unknown-linux-gnu/offload/lib:$LD_LIBRARY_PATH" + +RUSTFLAGS="-Ctarget-cpu=sm_90a -C target-feature=+ptx80 --emit=llvm-bc,llvm-ir -Zoffload=Device -Csave-temps -Zunstable-options" \ +cargo +offload build -Zunstable-options -r -v --target nvptx64-nvidia-cuda -Zbuild-std=core,panic_abort "${FEATURES_ARGS[@]}" + +RUSTFLAGS="--emit=llvm-bc,llvm-ir -Csave-temps -Zoffload=Host=/p/lustre1/drehwald1/prog/offload/rust_perf/target/nvptx64-nvidia-cuda/release/deps/device.bin -Zunstable-options \ +-Clink-arg=-L/tmp/drehwald1/prog/rust/build/x86_64-unknown-linux-gnu/offload/lib \ +-Clink-arg=-L/tmp/drehwald1/prog/rust/build/x86_64-unknown-linux-gnu/llvm/lib \ +-Clink-arg=-Wl,-rpath,/tmp/drehwald1/prog/rust/build/x86_64-unknown-linux-gnu/offload/lib \ +-Clink-arg=-Wl,-rpath,/tmp/drehwald1/prog/rust/build/x86_64-unknown-linux-gnu/llvm/lib \ +-Clink-arg=-lomptarget \ +-Clink-arg=-lomp" \ +cargo +offload build -r -Zbuild-std=std,panic_abort "${FEATURES_ARGS[@]}" + + +$CLANG_LINKER_WRAPPER "--should-extract=sm_90a" "--device-compiler=nvptx64-nvidia-cuda=-save-temps=cwd" "--device-compiler=nvptx-nvidia-cuda=-v" "--device-compiler=nvptx64-nvidia-cuda=-v" "${PTX80_ARGS[@]}" "--device-linker=nvptx-nvidia-cuda=-lompdevice" "--device-linker=nvptx-nvidia-cuda=--feature=+ptx80" "--device-compiler=nvptx64-nvidia-cuda=-march=sm_90a" "--device-compiler=nvptx64-nvidia-cuda=-Xclang" "--device-compiler=nvptx64-nvidia-cuda=-target-feature" "--device-compiler=nvptx64-nvidia-cuda=-Xclang" "--device-compiler=nvptx64-nvidia-cuda=+ptx80" "--host-triple=x86_64-unknown-linux-gnu" "--save-temps" "--linker-path=/tmp/drehwald1/prog/rust/build/x86_64-unknown-linux-gnu/lld/bin/ld.lld" "--hash-style=gnu" "--eh-frame-hdr" "-m" "elf_x86_64" "-pie" "-dynamic-linker" "/lib64/ld-linux-x86-64.so.2" "-o" "bare" "/lib/../lib64/Scrt1.o" "/lib/../lib64/crti.o" "/opt/rh/gcc-toolset-12/root/usr/lib/gcc/x86_64-redhat-linux/12/crtbeginS.o" "-L/tmp/drehwald1/prog/rust/build/x86_64-unknown-linux-gnu/llvm/bin/../lib/x86_64-unknown-linux-gnu" "-L/tmp/drehwald1/prog/rust/build/x86_64-unknown-linux-gnu/llvm/lib/clang/22/lib/x86_64-unknown-linux-gnu" "-L/opt/rh/gcc-toolset-12/root/usr/lib/gcc/x86_64-redhat-linux/12" "-L/opt/rh/gcc-toolset-12/root/usr/lib/gcc/x86_64-redhat-linux/12/../../../../lib64" "-L/lib/../lib64" "-L/usr/lib64" "-L/lib" "-L/usr/lib" "./target/release/deps/host.o" "-lstdc++" "-lm" "-L/tmp/drehwald1/prog/rust/build/x86_64-unknown-linux-gnu/offload/lib" "-lomp" "-lomptarget" "-L/tmp/drehwald1/prog/rust/build/x86_64-unknown-linux-gnu/llvm/lib" "-lgcc_s" "-lgcc" "-lpthread" "-lc" "-lgcc_s" "-lgcc" "/opt/rh/gcc-toolset-12/root/usr/lib/gcc/x86_64-redhat-linux/12/crtendS.o" "/lib/../lib64/crtn.o" + diff --git a/scripts/steps_dec.txt b/scripts/steps_dec.txt new file mode 100755 index 0000000..b3a32c9 --- /dev/null +++ b/scripts/steps_dec.txt @@ -0,0 +1,8 @@ +set -e +RUSTFLAGS="-Ctarget-cpu=gfx90a --emit=llvm-bc,llvm-ir -Zoffload=Device -Csave-temps -Zunstable-options" cargo +offload build -Zunstable-options -v --target amdgcn-amd-amdhsa -Zbuild-std=core -r + +RUSTFLAGS="--emit=llvm-bc,llvm-ir -Csave-temps -Zoffload=Host=/p/lustre1/drehwald1/prog/offload/rust_memtransfer/target/amdgcn-amd-amdhsa/release/deps/device.bin -Zunstable-options" cargo +offload build -r + +"/tmp/drehwald1/prog/rust/build/x86_64-unknown-linux-gnu/llvm/bin/clang-linker-wrapper" "--should-extract=gfx90a" "--device-compiler=amdgcn-amd-amdhsa=-g" "--device-compiler=amdgcn-amd-amdhsa=-save-temps=cwd" "--host-triple=x86_64-unknown-linux-gnu" "--save-temps" "--linker-path=/tmp/drehwald1/prog/rust/build/x86_64-unknown-linux-gnu/lld/bin/ld.lld" "--hash-style=gnu" "--eh-frame-hdr" "-m" "elf_x86_64" "-pie" "-dynamic-linker" "/lib64/ld-linux-x86-64.so.2" "-o" "bare" "/lib/../lib64/Scrt1.o" "/lib/../lib64/crti.o" "/opt/rh/gcc-toolset-12/root/usr/lib/gcc/x86_64-redhat-linux/12/crtbeginS.o" "-L/tmp/drehwald1/prog/rust/build/x86_64-unknown-linux-gnu/llvm/bin/../lib/x86_64-unknown-linux-gnu" "-L/tmp/drehwald1/prog/rust/build/x86_64-unknown-linux-gnu/llvm/lib/clang/21/lib/x86_64-unknown-linux-gnu" "-L/opt/rh/gcc-toolset-12/root/usr/lib/gcc/x86_64-redhat-linux/12" "-L/opt/rh/gcc-toolset-12/root/usr/lib/gcc/x86_64-redhat-linux/12/../../../../lib64" "-L/lib/../lib64" "-L/usr/lib64" "-L/lib" "-L/usr/lib" "./target/release/deps/host.o" "-lstdc++" "-lm" "-L/tmp/drehwald1/prog/rust/build/x86_64-unknown-linux-gnu/offload/lib" "-lomp" "-lomptarget" "-L/tmp/drehwald1/prog/rust/build/x86_64-unknown-linux-gnu/llvm/lib" "-lgcc_s" "-lgcc" "-lpthread" "-lc" "-lgcc_s" "-lgcc" "/opt/rh/gcc-toolset-12/root/usr/lib/gcc/x86_64-redhat-linux/12/crtendS.o" "/lib/../lib64/crtn.o" + +LIBOMPTARGET_INFO=-1 OFFLOAD_TRACK_ALLOCATION_TRACES=true ./bare diff --git a/src/apps/del_dot_vec_2d.rs b/src/apps/del_dot_vec_2d.rs index 6add513..9833157 100644 --- a/src/apps/del_dot_vec_2d.rs +++ b/src/apps/del_dot_vec_2d.rs @@ -17,43 +17,64 @@ const N_REAL_ZONES: usize = (IMAX - IMIN) * (JMAX - JMIN); const THREADS_PER_BLOCK: u32 = 256; const BLOCKS: u32 = (N_REAL_ZONES as u32).div_ceil(THREADS_PER_BLOCK); +use core::offload::offload_kernel; + +#[cfg(target_os = "linux")] +use rustc_offload_frontend::offload; + +#[cfg(target_os = "linux")] +use core::offload::offload::{preload, preload_mut, Preload, PreloadMut}; + +#[cfg(target_arch = "amdgpu")] +#[inline(always)] +fn block_dim_x() -> u32 { + use rustc_offload_frontend::gpu::*; + let dispatch = dispatch_ptr(); + let x = (*dispatch).workgroup_size_x as u32; + x +} + #[cfg(target_arch = "amdgpu")] use core::arch::amdgpu::{workgroup_id_x as block_idx_x, workitem_id_x as thread_idx_x}; + #[cfg(target_arch = "nvptx64")] use core::arch::nvptx::{ _block_dim_x as block_dim_x, _block_idx_x as block_idx_x, _thread_idx_x as thread_idx_x, }; -#[cfg(target_arch = "amdgpu")] -#[allow(improper_ctypes)] -unsafe extern "C" { - #[link_name = "llvm.amdgcn.workgroup.size.x"] - fn block_dim_x() -> u32; -} #[cfg(target_os = "linux")] use crate::common::data_utils::{ alloc, alloc_and_init_data, alloc_and_init_data_const, calc_checksum, free, inc_data_init_count, }; + #[cfg(target_os = "linux")] use crate::common::kernel_base::KernelBase; + #[cfg(target_os = "linux")] use crate::kernel_name; #[cfg(target_os = "linux")] -use crate::common::types::{Real, to_real}; +use crate::common::types::{to_real, Real}; #[cfg(target_os = "linux")] -pub struct DelDotVec2D { +pub struct DelDotVec2D<'a> { x: *mut Real, y: *mut Real, xdot: *mut Real, ydot: *mut Real, div: *mut Real, real_zones: *mut usize, + + p_x: Option>, + p_y: Option>, + p_xdot: Option>, + p_ydot: Option>, + p_real_zones: Option>, + p_div: Option>, } #[cfg(target_os = "linux")] -impl DelDotVec2D { +impl<'a> DelDotVec2D<'a> { pub const INIT: Self = DelDotVec2D { x: core::ptr::null_mut(), y: core::ptr::null_mut(), @@ -61,17 +82,26 @@ impl DelDotVec2D { ydot: core::ptr::null_mut(), div: core::ptr::null_mut(), real_zones: core::ptr::null_mut(), + + p_x: None, + p_y: None, + p_xdot: None, + p_ydot: None, + p_real_zones: None, + p_div: None, }; } #[cfg(target_os = "linux")] -impl KernelBase for DelDotVec2D { +impl<'a> KernelBase for DelDotVec2D<'a> { fn name(&self) -> &'static str { kernel_name!("DEL_DOT_VEC_2D") } + fn default_problem_size(&self) -> usize { N_REAL_ZONES } + fn default_reps(&self) -> u32 { DEFAULT_REPS } @@ -81,23 +111,29 @@ impl KernelBase for DelDotVec2D { self.x = alloc_and_init_data_const(NNALLS, to_real(0.0)); self.y = alloc_and_init_data_const(NNALLS, to_real(0.0)); } + self.real_zones = unsafe { alloc::(N_REAL_ZONES) }; + for i in 0..N_REAL_ZONES { unsafe { *self.real_zones.add(i) = usize::MAX; } } + inc_data_init_count(); let dx: f64 = 0.2; let dy: f64 = 0.1; + let jstart = JMIN as isize - NPNL as isize; let jend = JMAX as isize + 1 + NPNR as isize; let istart = IMIN as isize - NPNL as isize; let iend = IMAX as isize + 1 + NPNR as isize; + for j in jstart..jend { for i in istart..iend { let idx = (i + j * JP as isize) as usize; + unsafe { *self.x.add(idx) = to_real(i as f64 * dx); *self.y.add(idx) = to_real(j as f64 * dy); @@ -106,10 +142,12 @@ impl KernelBase for DelDotVec2D { } let j_stride = (IMAX - IMIN) as isize; + for j in JMIN..JMAX { for i in IMIN..IMAX { let iz = i + j * JP; let il = (i - IMIN) + (j - JMIN) * j_stride as usize; + unsafe { *self.real_zones.add(il) = iz; } @@ -121,61 +159,85 @@ impl KernelBase for DelDotVec2D { self.ydot = alloc_and_init_data(NNALLS); self.div = alloc_and_init_data_const(NNALLS, to_real(0.0)); } + + let x_ref: &'a [Real; NNALLS] = unsafe { &*(self.x as *const [Real; NNALLS]) }; + + let y_ref: &'a [Real; NNALLS] = unsafe { &*(self.y as *const [Real; NNALLS]) }; + + let xdot_ref: &'a [Real; NNALLS] = unsafe { &*(self.xdot as *const [Real; NNALLS]) }; + + let ydot_ref: &'a [Real; NNALLS] = unsafe { &*(self.ydot as *const [Real; NNALLS]) }; + + let real_zones_ref: &'a [usize; N_REAL_ZONES] = + unsafe { &*(self.real_zones as *const [usize; N_REAL_ZONES]) }; + + let div_ref: &'a mut [Real; NNALLS] = unsafe { &mut *(self.div as *mut [Real; NNALLS]) }; + + self.p_x = Some(preload(x_ref)); + self.p_y = Some(preload(y_ref)); + self.p_xdot = Some(preload(xdot_ref)); + self.p_ydot = Some(preload(ydot_ref)); + self.p_real_zones = Some(preload(real_zones_ref)); + self.p_div = Some(preload_mut(div_ref)); } fn run_kernel(&mut self) { let ptiny = to_real(1.0e-20); let half = to_real(0.5); - let x1 = self.x as *const [Real; NNALLS]; - let x2 = unsafe { self.x.add(1) as *const Real }; - let x3 = unsafe { self.x.add(1 + JP) as *const Real }; - let x4 = unsafe { self.x.add(JP) as *const Real }; - - let y1 = self.y as *const [Real; NNALLS]; - let y2 = unsafe { self.y.add(1) as *const Real }; - let y3 = unsafe { self.y.add(1 + JP) as *const Real }; - let y4 = unsafe { self.y.add(JP) as *const Real }; - - let fx1 = self.xdot as *const [Real; NNALLS]; - let fx2 = unsafe { self.xdot.add(1) as *const Real }; - let fx3 = unsafe { self.xdot.add(1 + JP) as *const Real }; - let fx4 = unsafe { self.xdot.add(JP) as *const Real }; - - let fy1 = self.ydot as *const [Real; NNALLS]; - let fy2 = unsafe { self.ydot.add(1) as *const Real }; - let fy3 = unsafe { self.ydot.add(1 + JP) as *const Real }; - let fy4 = unsafe { self.ydot.add(JP) as *const Real }; + let _p_x = self + .p_x + .as_ref() + .expect("DEL_DOT_VEC_2D x was not preloaded"); + + let _p_y = self + .p_y + .as_ref() + .expect("DEL_DOT_VEC_2D y was not preloaded"); + + let _p_xdot = self + .p_xdot + .as_ref() + .expect("DEL_DOT_VEC_2D xdot was not preloaded"); + + let _p_ydot = self + .p_ydot + .as_ref() + .expect("DEL_DOT_VEC_2D ydot was not preloaded"); + + let _p_real_zones = self + .p_real_zones + .as_ref() + .expect("DEL_DOT_VEC_2D real_zones was not preloaded"); + + let _p_div = self + .p_div + .as_ref() + .expect("DEL_DOT_VEC_2D div was not preloaded"); + + let x = unsafe { &*(self.x as *const [Real; NNALLS]) }; + let y = unsafe { &*(self.y as *const [Real; NNALLS]) }; + let xdot = unsafe { &*(self.xdot as *const [Real; NNALLS]) }; + let ydot = unsafe { &*(self.ydot as *const [Real; NNALLS]) }; + let real_zones = unsafe { &*(self.real_zones as *const [usize; N_REAL_ZONES]) }; unsafe { - core::intrinsics::offload::<_, _, ()>( - _del_dot_vec_2d, - [BLOCKS, 1, 1], - [THREADS_PER_BLOCK, 1, 1], - ( + offload! { + kernel = del_dot_vec_2d, + grid_dim = [BLOCKS, 1, 1], + block_dim = [THREADS_PER_BLOCK, 1, 1], + args = ( self.div as *mut [Real; NNALLS], - &*x1, - x2, - x3, - x4, - &*y1, - y2, - y3, - y4, - &*fx1, - fx2, - fx3, - fx4, - &*fy1, - fy2, - fy3, - fy4, - &*(self.real_zones as *const [usize; N_REAL_ZONES]), + x, + y, + xdot, + ydot, + real_zones, half, ptiny, N_REAL_ZONES, ), - ); + }; } } @@ -183,105 +245,121 @@ impl KernelBase for DelDotVec2D { unsafe { calc_checksum(self.div as *const Real, NNALLS) } } - fn tear_down(&mut self) { + fn tear_down(&mut self) -> f64 { + // Drop output preload first so device writes are copied back before checksum. + drop(self.p_div.take()); + + let checksum = self.update_checksum(); + + // Drop read-only preloads before freeing their host allocations. + drop(self.p_x.take()); + drop(self.p_y.take()); + drop(self.p_xdot.take()); + drop(self.p_ydot.take()); + drop(self.p_real_zones.take()); + unsafe { - free(self.x); - self.x = core::ptr::null_mut(); - free(self.y); - self.y = core::ptr::null_mut(); - free(self.xdot); - self.xdot = core::ptr::null_mut(); - free(self.ydot); - self.ydot = core::ptr::null_mut(); - free(self.div); - self.div = core::ptr::null_mut(); - free(self.real_zones); - self.real_zones = core::ptr::null_mut(); + if !self.x.is_null() { + free(self.x); + self.x = core::ptr::null_mut(); + } + + if !self.y.is_null() { + free(self.y); + self.y = core::ptr::null_mut(); + } + + if !self.xdot.is_null() { + free(self.xdot); + self.xdot = core::ptr::null_mut(); + } + + if !self.ydot.is_null() { + free(self.ydot); + self.ydot = core::ptr::null_mut(); + } + + if !self.div.is_null() { + free(self.div); + self.div = core::ptr::null_mut(); + } + + if !self.real_zones.is_null() { + free(self.real_zones); + self.real_zones = core::ptr::null_mut(); + } } - } -} -#[cfg(target_os = "linux")] -unsafe extern "C" { - pub fn _del_dot_vec_2d( - div: *mut [Real; NNALLS], - x1: &[Real; NNALLS], - x2: *const Real, - x3: *const Real, - x4: *const Real, - y1: &[Real; NNALLS], - y2: *const Real, - y3: *const Real, - y4: *const Real, - fx1: &[Real; NNALLS], - fx2: *const Real, - fx3: *const Real, - fx4: *const Real, - fy1: &[Real; NNALLS], - fy2: *const Real, - fy3: *const Real, - fy4: *const Real, - real_zones: &[usize; N_REAL_ZONES], - half: Real, - ptiny: Real, - iend: usize, - ); + checksum + } } #[cfg(not(target_os = "linux"))] use crate::common::types::{Real, RealExt}; -#[cfg(not(target_os = "linux"))] -#[unsafe(no_mangle)] -#[rustc_offload_kernel] -pub unsafe extern "gpu-kernel" fn _del_dot_vec_2d( +#[offload_kernel] +fn del_dot_vec_2d( div: *mut [Real; NNALLS], - x1: &[Real; NNALLS], - x2: *const Real, - x3: *const Real, - x4: *const Real, - y1: &[Real; NNALLS], - y2: *const Real, - y3: *const Real, - y4: *const Real, - fx1: &[Real; NNALLS], - fx2: *const Real, - fx3: *const Real, - fx4: *const Real, - fy1: &[Real; NNALLS], - fy2: *const Real, - fy3: *const Real, - fy4: *const Real, + x: &[Real; NNALLS], + y: &[Real; NNALLS], + xdot: &[Real; NNALLS], + ydot: &[Real; NNALLS], real_zones: &[usize; N_REAL_ZONES], half: Real, ptiny: Real, iend: usize, ) { let ii = unsafe { (block_idx_x() * block_dim_x() + thread_idx_x()) as usize }; + if ii < iend { + core::hint::assert_unchecked(ii < N_REAL_ZONES); let i = real_zones[ii]; - let xi = half * (x1[i] + *x2.add(i) - *x3.add(i) - *x4.add(i)); - let xj = half * (*x2.add(i) + *x3.add(i) - *x4.add(i) - x1[i]); + unsafe { + core::hint::assert_unchecked(i >= 0); + core::hint::assert_unchecked(i < NNALLS); + core::hint::assert_unchecked(i + 1 + JP < NNALLS); + let x1 = x[i]; + let x2 = x[i + 1]; + let x3 = x[i + 1 + JP]; + let x4 = x[i + JP]; + + let y1 = y[i]; + let y2 = y[i + 1]; + let y3 = y[i + 1 + JP]; + let y4 = y[i + JP]; - let yi = half * (y1[i] + *y2.add(i) - *y3.add(i) - *y4.add(i)); - let yj = half * (*y2.add(i) + *y3.add(i) - *y4.add(i) - y1[i]); + let fx1 = xdot[i]; + let fx2 = xdot[i + 1]; + let fx3 = xdot[i + 1 + JP]; + let fx4 = xdot[i + JP]; - let fxi = half * (fx1[i] + *fx2.add(i) - *fx3.add(i) - *fx4.add(i)); - let fxj = half * (*fx2.add(i) + *fx3.add(i) - *fx4.add(i) - fx1[i]); + let fy1 = ydot[i]; + let fy2 = ydot[i + 1]; + let fy3 = ydot[i + 1 + JP]; + let fy4 = ydot[i + JP]; - let fyi = half * (fy1[i] + *fy2.add(i) - *fy3.add(i) - *fy4.add(i)); - let fyj = half * (*fy2.add(i) + *fy3.add(i) - *fy4.add(i) - fy1[i]); + let xi = half * (x1 + x2 - x3 - x4); + let xj = half * (x2 + x3 - x4 - x1); - let rarea = Real::from(1.0) / (xi * yj - xj * yi + ptiny); + let yi = half * (y1 + y2 - y3 - y4); + let yj = half * (y2 + y3 - y4 - y1); - let dfxdx = rarea * (fxi * yj - fxj * yi); + let fxi = half * (fx1 + fx2 - fx3 - fx4); + let fxj = half * (fx2 + fx3 - fx4 - fx1); - let dfydy = rarea * (fyj * xi - fyi * xj); + let fyi = half * (fy1 + fy2 - fy3 - fy4); + let fyj = half * (fy2 + fy3 - fy4 - fy1); - let affine = (fy1[i] + *fy2.add(i) + *fy3.add(i) + *fy4.add(i)) - / (y1[i] + *y2.add(i) + *y3.add(i) + *y4.add(i)); + let rarea = Real::from(1.0) / (xi * yj - xj * yi + ptiny); - (*div)[i] = dfxdx + dfydy + affine; + let dfxdx = rarea * (fxi * yj - fxj * yi); + let dfydy = rarea * (fyj * xi - fyi * xj); + + let affine = (fy1 + fy2 + fy3 + fy4) / (y1 + y2 + y3 + y4); + + core::hint::assert_unchecked(i < NNALLS); + (*div)[i] = dfxdx + dfydy + affine; + } } } diff --git a/src/apps/energy.rs b/src/apps/energy.rs index 8fa9d51..bd564f4 100644 --- a/src/apps/energy.rs +++ b/src/apps/energy.rs @@ -1,5 +1,6 @@ #![allow(clippy::too_many_arguments)] #![cfg_attr(target_arch = "nvptx64", feature(stdarch_nvptx))] +#![cfg_attr(target_arch = "amdgpu", feature(stdarch_amdgpu))] const DEFAULT_PROBLEM_SIZE: usize = 1_000_000; const DEFAULT_REPS: u32 = 130; @@ -8,20 +9,14 @@ const IEND: usize = DEFAULT_PROBLEM_SIZE; const THREADS_PER_BLOCK: u32 = 256; const BLOCKS: u32 = (IEND as u32).div_ceil(THREADS_PER_BLOCK); -#[cfg(target_arch = "nvptx64")] -use core::arch::nvptx::{ - _block_dim_x as block_dim_x, _block_idx_x as block_idx_x, _thread_idx_x as thread_idx_x, -}; +use core::offload::offload_kernel; +use rustc_offload_frontend::partition::{Linear1D, PartitioningStrategy, Region}; -#[cfg(target_arch = "amdgpu")] -use core::arch::amdgpu::{workgroup_id_x as block_idx_x, workitem_id_x as thread_idx_x}; +#[cfg(target_os = "linux")] +use rustc_offload_frontend::offload; -#[cfg(target_arch = "amdgpu")] -#[allow(improper_ctypes)] -unsafe extern "C" { - #[link_name = "llvm.amdgcn.workgroup.size.x"] - fn block_dim_x() -> u32; -} +#[cfg(target_os = "linux")] +use core::offload::offload::{preload, preload_mut, Preload, PreloadMut}; #[cfg(target_os = "linux")] use crate::common::data_utils::{ @@ -33,14 +28,15 @@ use crate::common::kernel_base::KernelBase; use crate::kernel_name; #[cfg(target_os = "linux")] -use crate::common::types::{Real, to_real}; +use crate::common::types::{to_real, Real}; #[cfg(target_os = "linux")] -pub struct Energy { +pub struct Energy<'a> { rho0: Real, e_cut: Real, emin: Real, q_cut: Real, + e_new: *mut Real, e_old: *mut Real, delvc: *mut Real, @@ -56,15 +52,33 @@ pub struct Energy { ql_old: *mut Real, qq_old: *mut Real, vnewc: *mut Real, + + p_e_new: Option>, + p_q_new: Option>, + + p_e_old: Option>, + p_delvc: Option>, + p_p_new: Option>, + p_p_old: Option>, + p_q_old: Option>, + p_work: Option>, + p_comp_half_step: Option>, + p_p_half_step: Option>, + p_bvc: Option>, + p_pbvc: Option>, + p_ql_old: Option>, + p_qq_old: Option>, + p_vnewc: Option>, } #[cfg(target_os = "linux")] -impl Energy { +impl<'a> Energy<'a> { pub const INIT: Self = Energy { rho0: to_real(0.0), e_cut: to_real(0.0), emin: to_real(0.0), q_cut: to_real(0.0), + e_new: core::ptr::null_mut(), e_old: core::ptr::null_mut(), delvc: core::ptr::null_mut(), @@ -80,17 +94,36 @@ impl Energy { ql_old: core::ptr::null_mut(), qq_old: core::ptr::null_mut(), vnewc: core::ptr::null_mut(), + + p_e_new: None, + p_q_new: None, + + p_e_old: None, + p_delvc: None, + p_p_new: None, + p_p_old: None, + p_q_old: None, + p_work: None, + p_comp_half_step: None, + p_p_half_step: None, + p_bvc: None, + p_pbvc: None, + p_ql_old: None, + p_qq_old: None, + p_vnewc: None, }; } #[cfg(target_os = "linux")] -impl KernelBase for Energy { +impl<'a> KernelBase for Energy<'a> { fn name(&self) -> &'static str { kernel_name!("ENERGY") } + fn default_problem_size(&self) -> usize { DEFAULT_PROBLEM_SIZE } + fn default_reps(&self) -> u32 { DEFAULT_REPS } @@ -123,80 +156,215 @@ impl KernelBase for Energy { self.emin = init_data_scalar(); self.q_cut = init_data_scalar(); } + + self.p_e_new = Some(preload_mut(unsafe { + &mut *(self.e_new as *mut [Real; IEND]) + })); + self.p_q_new = Some(preload_mut(unsafe { + &mut *(self.q_new as *mut [Real; IEND]) + })); + + self.p_e_old = Some(preload(unsafe { + &*(self.e_old as *const [Real; IEND]) + })); + self.p_delvc = Some(preload(unsafe { + &*(self.delvc as *const [Real; IEND]) + })); + self.p_p_new = Some(preload(unsafe { + &*(self.p_new as *const [Real; IEND]) + })); + self.p_p_old = Some(preload(unsafe { + &*(self.p_old as *const [Real; IEND]) + })); + self.p_q_old = Some(preload(unsafe { + &*(self.q_old as *const [Real; IEND]) + })); + self.p_work = Some(preload(unsafe { + &*(self.work as *const [Real; IEND]) + })); + self.p_comp_half_step = Some(preload(unsafe { + &*(self.comp_half_step as *const [Real; IEND]) + })); + self.p_p_half_step = Some(preload(unsafe { + &*(self.p_half_step as *const [Real; IEND]) + })); + self.p_bvc = Some(preload(unsafe { + &*(self.bvc as *const [Real; IEND]) + })); + self.p_pbvc = Some(preload(unsafe { + &*(self.pbvc as *const [Real; IEND]) + })); + self.p_ql_old = Some(preload(unsafe { + &*(self.ql_old as *const [Real; IEND]) + })); + self.p_qq_old = Some(preload(unsafe { + &*(self.qq_old as *const [Real; IEND]) + })); + self.p_vnewc = Some(preload(unsafe { + &*(self.vnewc as *const [Real; IEND]) + })); } fn run_kernel(&mut self) { + let p_e_new = self + .p_e_new + .as_ref() + .expect("ENERGY e_new was not preloaded"); + + let p_q_new = self + .p_q_new + .as_ref() + .expect("ENERGY q_new was not preloaded"); + + self.p_e_old + .as_ref() + .expect("ENERGY e_old was not preloaded"); + self.p_delvc + .as_ref() + .expect("ENERGY delvc was not preloaded"); + self.p_p_new + .as_ref() + .expect("ENERGY p_new was not preloaded"); + self.p_p_old + .as_ref() + .expect("ENERGY p_old was not preloaded"); + self.p_q_old + .as_ref() + .expect("ENERGY q_old was not preloaded"); + self.p_work + .as_ref() + .expect("ENERGY work was not preloaded"); + self.p_comp_half_step + .as_ref() + .expect("ENERGY comp_half_step was not preloaded"); + self.p_p_half_step + .as_ref() + .expect("ENERGY p_half_step was not preloaded"); + self.p_bvc + .as_ref() + .expect("ENERGY bvc was not preloaded"); + self.p_pbvc + .as_ref() + .expect("ENERGY pbvc was not preloaded"); + self.p_ql_old + .as_ref() + .expect("ENERGY ql_old was not preloaded"); + self.p_qq_old + .as_ref() + .expect("ENERGY qq_old was not preloaded"); + self.p_vnewc + .as_ref() + .expect("ENERGY vnewc was not preloaded"); + + let mut e_new_reg = Region::<'_, _, Linear1D>::from(p_e_new); + let mut q_new_reg = Region::<'_, _, Linear1D>::from(p_q_new); + unsafe { - energycalc1( - self.e_new as *mut [Real; IEND], - &*(self.e_old as *const [Real; IEND]), - &*(self.delvc as *const [Real; IEND]), - &*(self.p_old as *const [Real; IEND]), - &*(self.q_old as *const [Real; IEND]), - &*(self.work as *const [Real; IEND]), - IEND, - ); - energycalc2( - &*(self.delvc as *const [Real; IEND]), - self.q_new as *mut [Real; IEND], - &*(self.comp_half_step as *const [Real; IEND]), - &*(self.p_half_step as *const [Real; IEND]), - self.e_new as *mut [Real; IEND], - &*(self.bvc as *const [Real; IEND]), - &*(self.pbvc as *const [Real; IEND]), - &*(self.ql_old as *const [Real; IEND]), - &*(self.qq_old as *const [Real; IEND]), - self.rho0, - IEND, - ); - energycalc3( - self.e_new as *mut [Real; IEND], - &*(self.delvc as *const [Real; IEND]), - &*(self.p_old as *const [Real; IEND]), - &*(self.q_old as *const [Real; IEND]), - &*(self.p_half_step as *const [Real; IEND]), - &*(self.q_new as *const [Real; IEND]), - IEND, - ); - energycalc4( - self.e_new as *mut [Real; IEND], - &*(self.work as *const [Real; IEND]), - self.e_cut, - self.emin, - IEND, - ); - energycalc5( - &*(self.delvc as *const [Real; IEND]), - &*(self.pbvc as *const [Real; IEND]), - self.e_new as *mut [Real; IEND], - &*(self.vnewc as *const [Real; IEND]), - &*(self.bvc as *const [Real; IEND]), - &*(self.p_new as *const [Real; IEND]), - &*(self.ql_old as *const [Real; IEND]), - &*(self.qq_old as *const [Real; IEND]), - &*(self.p_old as *const [Real; IEND]), - &*(self.q_old as *const [Real; IEND]), - &*(self.p_half_step as *const [Real; IEND]), - &*(self.q_new as *const [Real; IEND]), - self.rho0, - self.e_cut, - self.emin, - IEND, - ); - energycalc6( - &*(self.delvc as *const [Real; IEND]), - &*(self.pbvc as *const [Real; IEND]), - self.e_new as *mut [Real; IEND], - &*(self.vnewc as *const [Real; IEND]), - &*(self.bvc as *const [Real; IEND]), - &*(self.p_new as *const [Real; IEND]), - self.q_new as *mut [Real; IEND], - &*(self.ql_old as *const [Real; IEND]), - &*(self.qq_old as *const [Real; IEND]), - self.rho0, - self.q_cut, - IEND, - ); + offload! { + kernel = energycalc1, + grid_dim = [BLOCKS, 1, 1], + block_dim = [THREADS_PER_BLOCK, 1, 1], + args = ( + e_new_reg, + &*(self.e_old as *const [Real; IEND]), + &*(self.delvc as *const [Real; IEND]), + &*(self.p_old as *const [Real; IEND]), + &*(self.q_old as *const [Real; IEND]), + &*(self.work as *const [Real; IEND]), + IEND, + ), + }; + + offload! { + kernel = energycalc2, + grid_dim = [BLOCKS, 1, 1], + block_dim = [THREADS_PER_BLOCK, 1, 1], + args = ( + &*(self.delvc as *const [Real; IEND]), + q_new_reg, + &*(self.comp_half_step as *const [Real; IEND]), + &*(self.p_half_step as *const [Real; IEND]), + e_new_reg, + &*(self.bvc as *const [Real; IEND]), + &*(self.pbvc as *const [Real; IEND]), + &*(self.ql_old as *const [Real; IEND]), + &*(self.qq_old as *const [Real; IEND]), + self.rho0, + IEND, + ), + }; + + offload! { + kernel = energycalc3, + grid_dim = [BLOCKS, 1, 1], + block_dim = [THREADS_PER_BLOCK, 1, 1], + args = ( + e_new_reg, + &*(self.delvc as *const [Real; IEND]), + &*(self.p_old as *const [Real; IEND]), + &*(self.q_old as *const [Real; IEND]), + &*(self.p_half_step as *const [Real; IEND]), + &*(self.q_new as *const [Real; IEND]), + IEND, + ), + }; + + offload! { + kernel = energycalc4, + grid_dim = [BLOCKS, 1, 1], + block_dim = [THREADS_PER_BLOCK, 1, 1], + args = ( + e_new_reg, + &*(self.work as *const [Real; IEND]), + self.e_cut, + self.emin, + IEND, + ), + }; + + offload! { + kernel = energycalc5, + grid_dim = [BLOCKS, 1, 1], + block_dim = [THREADS_PER_BLOCK, 1, 1], + args = ( + &*(self.delvc as *const [Real; IEND]), + &*(self.pbvc as *const [Real; IEND]), + e_new_reg, + &*(self.vnewc as *const [Real; IEND]), + &*(self.bvc as *const [Real; IEND]), + &*(self.p_new as *const [Real; IEND]), + &*(self.ql_old as *const [Real; IEND]), + &*(self.qq_old as *const [Real; IEND]), + &*(self.p_old as *const [Real; IEND]), + &*(self.q_old as *const [Real; IEND]), + &*(self.p_half_step as *const [Real; IEND]), + &*(self.q_new as *const [Real; IEND]), + self.rho0, + self.e_cut, + self.emin, + IEND, + ), + }; + + offload! { + kernel = energycalc6, + grid_dim = [BLOCKS, 1, 1], + block_dim = [THREADS_PER_BLOCK, 1, 1], + args = ( + &*(self.delvc as *const [Real; IEND]), + &*(self.pbvc as *const [Real; IEND]), + e_new_reg, + &*(self.vnewc as *const [Real; IEND]), + &*(self.bvc as *const [Real; IEND]), + &*(self.p_new as *const [Real; IEND]), + q_new_reg, + &*(self.ql_old as *const [Real; IEND]), + &*(self.qq_old as *const [Real; IEND]), + self.rho0, + self.q_cut, + IEND, + ), + }; } } @@ -207,275 +375,112 @@ impl KernelBase for Energy { } } - fn tear_down(&mut self) { + fn tear_down(&mut self) -> f64 { + drop(self.p_e_new.take()); + drop(self.p_q_new.take()); + + drop(self.p_e_old.take()); + drop(self.p_delvc.take()); + drop(self.p_p_new.take()); + drop(self.p_p_old.take()); + drop(self.p_q_old.take()); + drop(self.p_work.take()); + drop(self.p_comp_half_step.take()); + drop(self.p_p_half_step.take()); + drop(self.p_bvc.take()); + drop(self.p_pbvc.take()); + drop(self.p_ql_old.take()); + drop(self.p_qq_old.take()); + drop(self.p_vnewc.take()); + + let ck = self.update_checksum(); + unsafe { - free(self.e_new); - self.e_new = core::ptr::null_mut(); - free(self.e_old); - self.e_old = core::ptr::null_mut(); - free(self.delvc); - self.delvc = core::ptr::null_mut(); - free(self.p_new); - self.p_new = core::ptr::null_mut(); - free(self.p_old); - self.p_old = core::ptr::null_mut(); - free(self.q_new); - self.q_new = core::ptr::null_mut(); - free(self.q_old); - self.q_old = core::ptr::null_mut(); - free(self.work); - self.work = core::ptr::null_mut(); - free(self.comp_half_step); - self.comp_half_step = core::ptr::null_mut(); - free(self.p_half_step); - self.p_half_step = core::ptr::null_mut(); - free(self.bvc); - self.bvc = core::ptr::null_mut(); - free(self.pbvc); - self.pbvc = core::ptr::null_mut(); - free(self.ql_old); - self.ql_old = core::ptr::null_mut(); - free(self.qq_old); - self.qq_old = core::ptr::null_mut(); - free(self.vnewc); - self.vnewc = core::ptr::null_mut(); - } - } -} + if !self.e_new.is_null() { + free(self.e_new); + self.e_new = core::ptr::null_mut(); + } -#[cfg(target_os = "linux")] -unsafe fn energycalc1( - e_new: *mut [Real; IEND], - e_old: &[Real; IEND], - delvc: &[Real; IEND], - p_old: &[Real; IEND], - q_old: &[Real; IEND], - work: &[Real; IEND], - iend: usize, -) { - core::intrinsics::offload( - _energycalc1, - [BLOCKS, 1, 1], - [THREADS_PER_BLOCK, 1, 1], - (e_new, e_old, delvc, p_old, q_old, work, iend), - ) -} -#[cfg(target_os = "linux")] -unsafe fn energycalc2( - delvc: &[Real; IEND], - q_new: *mut [Real; IEND], - comp_half_step: &[Real; IEND], - p_half_step: &[Real; IEND], - e_new: *mut [Real; IEND], - bvc: &[Real; IEND], - pbvc: &[Real; IEND], - ql_old: &[Real; IEND], - qq_old: &[Real; IEND], - rho0: Real, - iend: usize, -) { - core::intrinsics::offload( - _energycalc2, - [BLOCKS, 1, 1], - [THREADS_PER_BLOCK, 1, 1], - ( - delvc, - q_new, - comp_half_step, - p_half_step, - e_new, - bvc, - pbvc, - ql_old, - qq_old, - rho0, - iend, - ), - ) -} -#[cfg(target_os = "linux")] -unsafe fn energycalc3( - e_new: *mut [Real; IEND], - delvc: &[Real; IEND], - p_old: &[Real; IEND], - q_old: &[Real; IEND], - p_half_step: &[Real; IEND], - q_new: &[Real; IEND], - iend: usize, -) { - core::intrinsics::offload( - _energycalc3, - [BLOCKS, 1, 1], - [THREADS_PER_BLOCK, 1, 1], - (e_new, delvc, p_old, q_old, p_half_step, q_new, iend), - ) -} -#[cfg(target_os = "linux")] -unsafe fn energycalc4( - e_new: *mut [Real; IEND], - work: &[Real; IEND], - e_cut: Real, - emin: Real, - iend: usize, -) { - core::intrinsics::offload( - _energycalc4, - [BLOCKS, 1, 1], - [THREADS_PER_BLOCK, 1, 1], - (e_new, work, e_cut, emin, iend), - ) -} -#[cfg(target_os = "linux")] -unsafe fn energycalc5( - delvc: &[Real; IEND], - pbvc: &[Real; IEND], - e_new: *mut [Real; IEND], - vnewc: &[Real; IEND], - bvc: &[Real; IEND], - p_new: &[Real; IEND], - ql_old: &[Real; IEND], - qq_old: &[Real; IEND], - p_old: &[Real; IEND], - q_old: &[Real; IEND], - p_half_step: &[Real; IEND], - q_new: &[Real; IEND], - rho0: Real, - e_cut: Real, - emin: Real, - iend: usize, -) { - core::intrinsics::offload( - _energycalc5, - [BLOCKS, 1, 1], - [THREADS_PER_BLOCK, 1, 1], - ( - delvc, - pbvc, - e_new, - vnewc, - bvc, - p_new, - ql_old, - qq_old, - p_old, - q_old, - p_half_step, - q_new, - rho0, - e_cut, - emin, - iend, - ), - ) -} -#[cfg(target_os = "linux")] -unsafe fn energycalc6( - delvc: &[Real; IEND], - pbvc: &[Real; IEND], - e_new: *mut [Real; IEND], - vnewc: &[Real; IEND], - bvc: &[Real; IEND], - p_new: &[Real; IEND], - q_new: *mut [Real; IEND], - ql_old: &[Real; IEND], - qq_old: &[Real; IEND], - rho0: Real, - q_cut: Real, - iend: usize, -) { - core::intrinsics::offload( - _energycalc6, - [BLOCKS, 1, 1], - [THREADS_PER_BLOCK, 1, 1], - ( - delvc, pbvc, e_new, vnewc, bvc, p_new, q_new, ql_old, qq_old, rho0, q_cut, iend, - ), - ) -} + if !self.e_old.is_null() { + free(self.e_old); + self.e_old = core::ptr::null_mut(); + } -#[cfg(target_os = "linux")] -unsafe extern "C" { - pub fn _energycalc1( - e_new: *mut [Real; IEND], - e_old: &[Real; IEND], - delvc: &[Real; IEND], - p_old: &[Real; IEND], - q_old: &[Real; IEND], - work: &[Real; IEND], - iend: usize, - ); - pub fn _energycalc2( - delvc: &[Real; IEND], - q_new: *mut [Real; IEND], - comp_half_step: &[Real; IEND], - p_half_step: &[Real; IEND], - e_new: *mut [Real; IEND], - bvc: &[Real; IEND], - pbvc: &[Real; IEND], - ql_old: &[Real; IEND], - qq_old: &[Real; IEND], - rho0: Real, - iend: usize, - ); - pub fn _energycalc3( - e_new: *mut [Real; IEND], - delvc: &[Real; IEND], - p_old: &[Real; IEND], - q_old: &[Real; IEND], - p_half_step: &[Real; IEND], - q_new: &[Real; IEND], - iend: usize, - ); - pub fn _energycalc4( - e_new: *mut [Real; IEND], - work: &[Real; IEND], - e_cut: Real, - emin: Real, - iend: usize, - ); - pub fn _energycalc5( - delvc: &[Real; IEND], - pbvc: &[Real; IEND], - e_new: *mut [Real; IEND], - vnewc: &[Real; IEND], - bvc: &[Real; IEND], - p_new: &[Real; IEND], - ql_old: &[Real; IEND], - qq_old: &[Real; IEND], - p_old: &[Real; IEND], - q_old: &[Real; IEND], - p_half_step: &[Real; IEND], - q_new: &[Real; IEND], - rho0: Real, - e_cut: Real, - emin: Real, - iend: usize, - ); - pub fn _energycalc6( - delvc: &[Real; IEND], - pbvc: &[Real; IEND], - e_new: *mut [Real; IEND], - vnewc: &[Real; IEND], - bvc: &[Real; IEND], - p_new: &[Real; IEND], - q_new: *mut [Real; IEND], - ql_old: &[Real; IEND], - qq_old: &[Real; IEND], - rho0: Real, - q_cut: Real, - iend: usize, - ); + if !self.delvc.is_null() { + free(self.delvc); + self.delvc = core::ptr::null_mut(); + } + + if !self.p_new.is_null() { + free(self.p_new); + self.p_new = core::ptr::null_mut(); + } + + if !self.p_old.is_null() { + free(self.p_old); + self.p_old = core::ptr::null_mut(); + } + + if !self.q_new.is_null() { + free(self.q_new); + self.q_new = core::ptr::null_mut(); + } + + if !self.q_old.is_null() { + free(self.q_old); + self.q_old = core::ptr::null_mut(); + } + + if !self.work.is_null() { + free(self.work); + self.work = core::ptr::null_mut(); + } + + if !self.comp_half_step.is_null() { + free(self.comp_half_step); + self.comp_half_step = core::ptr::null_mut(); + } + + if !self.p_half_step.is_null() { + free(self.p_half_step); + self.p_half_step = core::ptr::null_mut(); + } + + if !self.bvc.is_null() { + free(self.bvc); + self.bvc = core::ptr::null_mut(); + } + + if !self.pbvc.is_null() { + free(self.pbvc); + self.pbvc = core::ptr::null_mut(); + } + + if !self.ql_old.is_null() { + free(self.ql_old); + self.ql_old = core::ptr::null_mut(); + } + + if !self.qq_old.is_null() { + free(self.qq_old); + self.qq_old = core::ptr::null_mut(); + } + + if !self.vnewc.is_null() { + free(self.vnewc); + self.vnewc = core::ptr::null_mut(); + } + } + ck + } } #[cfg(not(target_os = "linux"))] use crate::common::types::{Real, RealExt}; -#[cfg(not(target_os = "linux"))] -#[unsafe(no_mangle)] -#[inline(never)] -#[rustc_offload_kernel] -pub extern "gpu-kernel" fn _energycalc1( - e_new: *mut [Real; IEND], +#[offload_kernel] +fn energycalc1( + mut e_new: Region, e_old: &[Real; IEND], delvc: &[Real; IEND], p_old: &[Real; IEND], @@ -483,25 +488,21 @@ pub extern "gpu-kernel" fn _energycalc1( work: &[Real; IEND], iend: usize, ) { - unsafe { - let i = (block_idx_x() * block_dim_x() + thread_idx_x()) as usize; - if i < iend { - (*e_new)[i] = (*e_old)[i] - Real::from(0.5) * (*delvc)[i] * ((*p_old)[i] + (*q_old)[i]) - + Real::from(0.5) * (*work)[i]; - } + let i = Linear1D::index(); + + if let Some(v) = e_new.get_mut() { + *v = (*e_old)[i] - Real::from(0.5) * (*delvc)[i] * ((*p_old)[i] + (*q_old)[i]) + + Real::from(0.5) * (*work)[i]; } } -#[cfg(not(target_os = "linux"))] -#[unsafe(no_mangle)] -#[inline(never)] -#[rustc_offload_kernel] -pub extern "gpu-kernel" fn _energycalc2( +#[offload_kernel] +fn energycalc2( delvc: &[Real; IEND], - q_new: *mut [Real; IEND], + mut q_new: Region, comp_half_step: &[Real; IEND], p_half_step: &[Real; IEND], - e_new: *mut [Real; IEND], + mut e_new: Region, bvc: &[Real; IEND], pbvc: &[Real; IEND], ql_old: &[Real; IEND], @@ -509,33 +510,32 @@ pub extern "gpu-kernel" fn _energycalc2( rho0: Real, iend: usize, ) { - unsafe { - let i = (block_idx_x() * block_dim_x() + thread_idx_x()) as usize; - if i < iend { - if ((*delvc)[i]).to_f64() > 0.0 { - (*q_new)[i] = Real::from(0.0); + let i = Linear1D::index(); + + if let Some(v1) = q_new.get_mut() + && let Some(v2) = e_new.get_mut() + { + if ((*delvc)[i]).to_f64() > 0.0 { + *v1 = Real::from(0.0); + } else { + let vhalf = Real::from(1.0) / (Real::from(1.0) + (*comp_half_step)[i]); + let mut ssc = + ((*pbvc)[i] * (*v2) + vhalf * vhalf * (*bvc)[i] * (*p_half_step)[i]) / rho0; + + if ssc.to_f64() <= 0.1111111e-36 { + ssc = Real::from(0.3333333e-18); } else { - let vhalf = Real::from(1.0) / (Real::from(1.0) + (*comp_half_step)[i]); - let mut ssc = ((*pbvc)[i] * (*e_new)[i] - + vhalf * vhalf * (*bvc)[i] * (*p_half_step)[i]) - / rho0; - if ssc.to_f64() <= 0.1111111e-36 { - ssc = Real::from(0.3333333e-18); - } else { - ssc = ssc.sqrt(); - } - (*q_new)[i] = ssc * (*ql_old)[i] + (*qq_old)[i]; + ssc = ssc.sqrt(); } + + *v1 = ssc * (*ql_old)[i] + (*qq_old)[i]; } } } -#[cfg(not(target_os = "linux"))] -#[unsafe(no_mangle)] -#[inline(never)] -#[rustc_offload_kernel] -pub extern "gpu-kernel" fn _energycalc3( - e_new: *mut [Real; IEND], +#[offload_kernel] +fn energycalc3( + mut e_new: Region, delvc: &[Real; IEND], p_old: &[Real; IEND], q_old: &[Real; IEND], @@ -543,50 +543,44 @@ pub extern "gpu-kernel" fn _energycalc3( q_new: &[Real; IEND], iend: usize, ) { - unsafe { - let i = (block_idx_x() * block_dim_x() + thread_idx_x()) as usize; - if i < iend { - (*e_new)[i] += Real::from(0.5) - * (*delvc)[i] - * (Real::from(3.0) * ((*p_old)[i] + (*q_old)[i]) - - Real::from(4.0) * ((*p_half_step)[i] + (*q_new)[i])); - } + let i = Linear1D::index(); + + if let Some(v) = e_new.get_mut() { + *v += Real::from(0.5) + * (*delvc)[i] + * (Real::from(3.0) * ((*p_old)[i] + (*q_old)[i]) + - Real::from(4.0) * ((*p_half_step)[i] + (*q_new)[i])); } } -#[cfg(not(target_os = "linux"))] -#[unsafe(no_mangle)] -#[inline(never)] -#[rustc_offload_kernel] -pub extern "gpu-kernel" fn _energycalc4( - e_new: *mut [Real; IEND], +#[offload_kernel] +fn energycalc4( + mut e_new: Region, work: &[Real; IEND], e_cut: Real, emin: Real, iend: usize, ) { - unsafe { - let i = (block_idx_x() * block_dim_x() + thread_idx_x()) as usize; - if i < iend { - (*e_new)[i] += Real::from(0.5) * (*work)[i]; - if ((*e_new)[i]).abs() < e_cut { - (*e_new)[i] = Real::from(0.0); - } - if (*e_new)[i] < emin { - (*e_new)[i] = emin; - } + let i = Linear1D::index(); + + if let Some(v) = e_new.get_mut() { + *v += Real::from(0.5) * (*work)[i]; + + if (*v).abs() < e_cut { + *v = Real::from(0.0); + } + + if *v < emin { + *v = emin; } } } -#[cfg(not(target_os = "linux"))] -#[unsafe(no_mangle)] -#[inline(never)] -#[rustc_offload_kernel] -pub extern "gpu-kernel" fn _energycalc5( +#[offload_kernel] +fn energycalc5( delvc: &[Real; IEND], pbvc: &[Real; IEND], - e_new: *mut [Real; IEND], + mut e_new: Region, vnewc: &[Real; IEND], bvc: &[Real; IEND], p_new: &[Real; IEND], @@ -601,70 +595,74 @@ pub extern "gpu-kernel" fn _energycalc5( emin: Real, iend: usize, ) { - unsafe { - let i = (block_idx_x() * block_dim_x() + thread_idx_x()) as usize; - if i < iend { - let q_tilde = if ((*delvc)[i]).to_f64() > 0.0 { - Real::from(0.0) + let i = Linear1D::index(); + + if let Some(v) = e_new.get_mut() { + let q_tilde = if ((*delvc)[i]).to_f64() > 0.0 { + Real::from(0.0) + } else { + let mut ssc = + ((*pbvc)[i] * (*v) + (*vnewc)[i] * (*vnewc)[i] * (*bvc)[i] * (*p_new)[i]) / rho0; + + if ssc.to_f64() <= 0.1111111e-36 { + ssc = Real::from(0.3333333e-18); } else { - let mut ssc = ((*pbvc)[i] * (*e_new)[i] - + (*vnewc)[i] * (*vnewc)[i] * (*bvc)[i] * (*p_new)[i]) - / rho0; - if ssc.to_f64() <= 0.1111111e-36 { - ssc = Real::from(0.3333333e-18); - } else { - ssc = ssc.sqrt(); - } - ssc * (*ql_old)[i] + (*qq_old)[i] - }; - (*e_new)[i] -= (Real::from(7.0) * ((*p_old)[i] + (*q_old)[i]) - - Real::from(8.0) * ((*p_half_step)[i] + (*q_new)[i]) - + ((*p_new)[i] + q_tilde)) - * (*delvc)[i] - / Real::from(6.0); - if ((*e_new)[i]).abs() < e_cut { - (*e_new)[i] = Real::from(0.0); - } - if (*e_new)[i] < emin { - (*e_new)[i] = emin; + ssc = ssc.sqrt(); } + + ssc * (*ql_old)[i] + (*qq_old)[i] + }; + + *v -= (Real::from(7.0) * ((*p_old)[i] + (*q_old)[i]) + - Real::from(8.0) * ((*p_half_step)[i] + (*q_new)[i]) + + ((*p_new)[i] + q_tilde)) + * (*delvc)[i] + / Real::from(6.0); + + if (*v).abs() < e_cut { + *v = Real::from(0.0); + } + + if *v < emin { + *v = emin; } } } -#[cfg(not(target_os = "linux"))] -#[unsafe(no_mangle)] -#[inline(never)] -#[rustc_offload_kernel] -pub extern "gpu-kernel" fn _energycalc6( +#[offload_kernel] +fn energycalc6( delvc: &[Real; IEND], pbvc: &[Real; IEND], - e_new: *mut [Real; IEND], + mut e_new: Region, vnewc: &[Real; IEND], bvc: &[Real; IEND], p_new: &[Real; IEND], - q_new: *mut [Real; IEND], + mut q_new: Region, ql_old: &[Real; IEND], qq_old: &[Real; IEND], rho0: Real, q_cut: Real, iend: usize, ) { - unsafe { - let i = (block_idx_x() * block_dim_x() + thread_idx_x()) as usize; - if i < iend && ((*delvc)[i]).to_f64() <= 0.0 { - let mut ssc = ((*pbvc)[i] * (*e_new)[i] - + (*vnewc)[i] * (*vnewc)[i] * (*bvc)[i] * (*p_new)[i]) - / rho0; - if ssc.to_f64() <= 0.1111111e-36 { - ssc = Real::from(0.3333333e-18); - } else { - ssc = ssc.sqrt(); - } - (*q_new)[i] = ssc * (*ql_old)[i] + (*qq_old)[i]; - if ((*q_new)[i]).abs() < q_cut { - (*q_new)[i] = Real::from(0.0); - } + let i = Linear1D::index(); + + if let Some(v1) = e_new.get_mut() + && let Some(v2) = q_new.get_mut() + && ((*delvc)[i]).to_f64() <= 0.0 + { + let mut ssc = + ((*pbvc)[i] * (*v1) + (*vnewc)[i] * (*vnewc)[i] * (*bvc)[i] * (*p_new)[i]) / rho0; + + if ssc.to_f64() <= 0.1111111e-36 { + ssc = Real::from(0.3333333e-18); + } else { + ssc = ssc.sqrt(); + } + + *v2 = ssc * (*ql_old)[i] + (*qq_old)[i]; + + if (*v2).abs() < q_cut { + *v2 = Real::from(0.0); } } } diff --git a/src/apps/fir.rs b/src/apps/fir.rs index 1b267b6..c9de91c 100644 --- a/src/apps/fir.rs +++ b/src/apps/fir.rs @@ -1,25 +1,22 @@ +#![feature(stmt_expr_attributes)] const DEFAULT_PROBLEM_SIZE: usize = 1_000_000; const DEFAULT_REPS: u32 = 160; const IEND: usize = DEFAULT_PROBLEM_SIZE; pub const COEFFLEN: usize = 16; +const INLEN: usize = IEND + COEFFLEN - 1; + const THREADS_PER_BLOCK: u32 = 256; const BLOCKS: u32 = (IEND as u32).div_ceil(THREADS_PER_BLOCK); -#[cfg(target_arch = "nvptx64")] -use core::arch::nvptx::{ - _block_dim_x as block_dim_x, _block_idx_x as block_idx_x, _thread_idx_x as thread_idx_x, -}; -#[cfg(target_arch = "amdgpu")] -#[allow(improper_ctypes)] -unsafe extern "C" { - #[link_name = "llvm.amdgcn.workitem.id.x"] - fn thread_idx_x() -> i32; - #[link_name = "llvm.amdgcn.workgroup.id.x"] - fn block_idx_x() -> i32; - #[link_name = "llvm.amdgcn.workgroup.size.x"] - fn block_dim_x() -> i32; -} +use core::offload::offload_kernel; +use rustc_offload_frontend::partition::{Linear1D, PartitioningStrategy, Region}; + +#[cfg(target_os = "linux")] +use rustc_offload_frontend::offload; + +#[cfg(target_os = "linux")] +use core::offload::offload::{preload, preload_mut, Preload, PreloadMut}; #[cfg(target_os = "linux")] use crate::common::data_utils::{ @@ -31,126 +28,180 @@ use crate::common::kernel_base::KernelBase; use crate::kernel_name; #[cfg(target_os = "linux")] -use crate::common::types::{Real, to_real}; +use crate::common::types::{to_real, Real}; #[cfg(target_os = "linux")] -pub struct Fir { +pub struct Fir<'a> { m_in: *mut Real, m_out: *mut Real, - coeff: [Real; COEFFLEN], + coeff: *mut Real, + + p_m_in: Option>, + p_m_out: Option>, + p_coeff: Option>, } #[cfg(target_os = "linux")] -impl Fir { +impl<'a> Fir<'a> { pub const INIT: Self = Fir { m_in: core::ptr::null_mut(), m_out: core::ptr::null_mut(), - coeff: [const { to_real(0.0) }; COEFFLEN], + coeff: core::ptr::null_mut(), + + p_m_in: None, + p_m_out: None, + p_coeff: None, }; } #[cfg(target_os = "linux")] -impl KernelBase for Fir { +impl<'a> KernelBase for Fir<'a> { fn name(&self) -> &'static str { kernel_name!("FIR") } + fn default_problem_size(&self) -> usize { DEFAULT_PROBLEM_SIZE } + fn default_reps(&self) -> u32 { DEFAULT_REPS } fn setup(&mut self) { - self.coeff = [ - to_real(3.0), - to_real(-1.0), - to_real(-1.0), - to_real(-1.0), - to_real(-1.0), - to_real(3.0), - to_real(-1.0), - to_real(-1.0), - to_real(-1.0), - to_real(-1.0), - to_real(3.0), - to_real(-1.0), - to_real(-1.0), - to_real(-1.0), - to_real(-1.0), - to_real(3.0), - ]; - unsafe { - self.m_in = alloc_and_init_data_rand_value(IEND + COEFFLEN - 1); + self.m_in = alloc_and_init_data_rand_value(INLEN); self.m_out = alloc_and_init_data_const(IEND, to_real(0.0)); + + // Allocate coeff on the heap as well, so the Preload handle does not + // borrow from inside `self`. + // + // This random-inits first, then we overwrite with the hard-coded FIR + // coefficients below. + self.coeff = alloc_and_init_data_rand_value(COEFFLEN); + + let coeff: &mut [Real; COEFFLEN] = &mut *(self.coeff as *mut [Real; COEFFLEN]); + + *coeff = [ + to_real(3.0), + to_real(-1.0), + to_real(-1.0), + to_real(-1.0), + to_real(-1.0), + to_real(3.0), + to_real(-1.0), + to_real(-1.0), + to_real(-1.0), + to_real(-1.0), + to_real(3.0), + to_real(-1.0), + to_real(-1.0), + to_real(-1.0), + to_real(-1.0), + to_real(3.0), + ]; } + + let m_in_ref: &'a [Real; INLEN] = unsafe { &*(self.m_in as *const [Real; INLEN]) }; + + let m_out_ref: &'a mut [Real; IEND] = unsafe { &mut *(self.m_out as *mut [Real; IEND]) }; + + let coeff_ref: &'a [Real; COEFFLEN] = unsafe { &*(self.coeff as *const [Real; COEFFLEN]) }; + + self.p_m_in = Some(preload(m_in_ref)); + self.p_m_out = Some(preload_mut(m_out_ref)); + self.p_coeff = Some(preload(coeff_ref)); } fn run_kernel(&mut self) { - unsafe { - core::intrinsics::offload::<_, _, ()>( - _fir, - [BLOCKS, 1, 1], - [THREADS_PER_BLOCK, 1, 1], - ( - self.m_out as *mut [Real; IEND], - &*(self.m_in as *const [Real; IEND + COEFFLEN]), - &self.coeff as &[Real; COEFFLEN], - IEND, - ), - ); - } + let Some(_p_m_in) = self.p_m_in.as_ref() else { + return; + }; + let Some(p_m_out) = self.p_m_out.as_ref() else { + return; + }; + let Some(_p_coeff) = self.p_coeff.as_ref() else { + return; + }; + + let mut m_out_reg = Region::<'_, _, Linear1D>::from(p_m_out); + + let m_in_ref: &[Real; INLEN] = unsafe { &*(self.m_in as *const [Real; INLEN]) }; + + let coeff_ref: &[Real; COEFFLEN] = unsafe { &*(self.coeff as *const [Real; COEFFLEN]) }; + + offload! { + kernel = fir, + grid_dim = [BLOCKS, 1, 1], + block_dim = [THREADS_PER_BLOCK, 1, 1], + args = ( + m_out_reg, + m_in_ref, + coeff_ref, + IEND, + ), + }; } fn update_checksum(&self) -> f64 { unsafe { calc_checksum(self.m_out as *const Real, IEND) } } - fn tear_down(&mut self) { + fn tear_down(&mut self) -> f64 { + // Drop all offload mappings before freeing the corresponding host memory. + drop(self.p_m_out.take()); + drop(self.p_m_in.take()); + drop(self.p_coeff.take()); + let ck = self.update_checksum(); + unsafe { - free(self.m_in); - self.m_in = core::ptr::null_mut(); - free(self.m_out); - self.m_out = core::ptr::null_mut(); + if !self.m_in.is_null() { + free(self.m_in); + self.m_in = core::ptr::null_mut(); + } + + if !self.m_out.is_null() { + free(self.m_out); + self.m_out = core::ptr::null_mut(); + } + + if !self.coeff.is_null() { + free(self.coeff); + self.coeff = core::ptr::null_mut(); + } } + ck } } -#[cfg(target_os = "linux")] -unsafe extern "C" { - pub fn _fir( - m_out: *mut [Real; IEND], - m_in: &[Real; IEND + COEFFLEN], - coeff: &[Real; COEFFLEN], - iend: usize, - ); -} - #[cfg(not(target_os = "linux"))] use crate::common::types::Real; -#[cfg(not(target_os = "linux"))] -#[unsafe(no_mangle)] -#[rustc_offload_kernel] -pub unsafe extern "gpu-kernel" fn _fir( - m_out: *mut [Real; IEND], - m_in: &[Real; IEND + COEFFLEN], +#[offload_kernel] +fn fir( + mut m_out: Region, + m_in: &[Real; INLEN], coeff: &[Real; COEFFLEN], iend: usize, ) { - let i = unsafe { (block_idx_x() * block_dim_x() + thread_idx_x()) as usize }; + let i = Linear1D::index(); + if i < iend { - let mut sum: Real = Real::from(0.0); - let mut j = 0; - while j < COEFFLEN { - unsafe { - sum += (*coeff)[j] * (*m_in)[i + j]; + if let Some(v) = m_out.get_mut() { + let mut sum: Real = Real::from(0.0); + let mut j = 0; + + //#[unroll(2)] + //for j in 0..COEFFLEN { + //#[rustc_unroll(4)] + while j < COEFFLEN { + unsafe { + sum += (*coeff)[j] * (*m_in)[i + j]; + } + j += 1; } - j += 1; - } - unsafe { - (*m_out)[i] = sum; + + *v = sum; } } } diff --git a/src/apps/ltimes.rs b/src/apps/ltimes.rs index deedc70..72caa34 100644 --- a/src/apps/ltimes.rs +++ b/src/apps/ltimes.rs @@ -1,73 +1,87 @@ pub const NUM_D: usize = 64; pub const NUM_G: usize = 32; pub const NUM_M: usize = 25; + +const DEFAULT_PROBLEM_SIZE: usize = 1_000_000; const DEFAULT_REPS: u32 = 50; +const DEFAULT_NUM_Z: usize = (DEFAULT_PROBLEM_SIZE + (NUM_D * NUM_G) / 2) / (NUM_D * NUM_G); + +const PHILEN: usize = NUM_M * NUM_G * DEFAULT_NUM_Z; +const ELLLEN: usize = NUM_D * NUM_M; +const PSILEN: usize = NUM_D * NUM_G * DEFAULT_NUM_Z; + +use core::offload::offload_kernel; +use rustc_offload_frontend::partition::{PartitioningStrategy, Region, Stride3D}; + +#[cfg(target_os = "linux")] +use rustc_offload_frontend::offload; + +#[cfg(target_os = "linux")] +use core::offload::offload::{preload, preload_mut, Preload, PreloadMut}; + +#[cfg(target_arch = "amdgpu")] +use core::arch::amdgpu::{ + workgroup_id_x as block_idx_x, workgroup_id_y as block_idx_y, workgroup_id_z as block_idx_z, + workitem_id_x as thread_idx_x, workitem_id_y as thread_idx_y, workitem_id_z as thread_idx_z, +}; + #[cfg(target_arch = "nvptx64")] use core::arch::nvptx::{ _block_idx_x as block_idx_x, _block_idx_y as block_idx_y, _block_idx_z as block_idx_z, _thread_idx_x as thread_idx_x, _thread_idx_y as thread_idx_y, _thread_idx_z as thread_idx_z, }; -#[cfg(target_arch = "amdgpu")] -#[allow(improper_ctypes)] -unsafe extern "C" { - #[link_name = "llvm.amdgcn.workgroup.id.x"] - fn block_idx_x() -> i32; - #[link_name = "llvm.amdgcn.workgroup.id.y"] - fn block_idx_y() -> i32; - #[link_name = "llvm.amdgcn.workgroup.id.z"] - fn block_idx_z() -> i32; - - #[link_name = "llvm.amdgcn.workitem.id.x"] - fn thread_idx_x() -> i32; - #[link_name = "llvm.amdgcn.workitem.id.y"] - fn thread_idx_y() -> i32; - #[link_name = "llvm.amdgcn.workitem.id.z"] - fn thread_idx_z() -> i32; -} - -#[cfg(target_os = "linux")] -use libc::printf; - #[cfg(target_os = "linux")] use crate::common::data_utils::{ alloc_and_init_data, alloc_and_init_data_const, calc_checksum, free, }; + #[cfg(target_os = "linux")] use crate::common::kernel_base::KernelBase; + #[cfg(target_os = "linux")] use crate::kernel_name; #[cfg(target_os = "linux")] -use crate::common::types::{Real, to_real}; +use crate::common::types::{to_real, Real}; #[cfg(target_os = "linux")] -pub struct LTimes { +pub struct LTimes<'a> { num_z: usize, + phidat: *mut Real, elldat: *mut Real, psidat: *mut Real, + + p_phidat: Option>, + p_elldat: Option>, + p_psidat: Option>, } #[cfg(target_os = "linux")] -impl LTimes { +impl<'a> LTimes<'a> { pub const INIT: Self = LTimes { num_z: 0, + phidat: core::ptr::null_mut(), elldat: core::ptr::null_mut(), psidat: core::ptr::null_mut(), + + p_phidat: None, + p_elldat: None, + p_psidat: None, }; } #[cfg(target_os = "linux")] -impl KernelBase for LTimes { +impl<'a> KernelBase for LTimes<'a> { fn name(&self) -> &'static str { kernel_name!("LTIMES") } fn default_problem_size(&self) -> usize { - let num_z_default = (1_000_000 + (NUM_D * NUM_G) / 2) / (NUM_D * NUM_G); + let num_z_default = (DEFAULT_PROBLEM_SIZE + (NUM_D * NUM_G) / 2) / (NUM_D * NUM_G); num_z_default * NUM_D * NUM_G } @@ -77,17 +91,50 @@ impl KernelBase for LTimes { fn setup(&mut self) { let prob_size = self.default_problem_size(); + self.num_z = (prob_size + (NUM_D * NUM_G) / 2) / (NUM_D * NUM_G); let philen = NUM_M * NUM_G * self.num_z; let elllen = NUM_D * NUM_M; let psilen = NUM_D * NUM_G * self.num_z; + assert_eq!( + self.num_z, DEFAULT_NUM_Z, + "LTIMES const-array preload version requires num_z = {DEFAULT_NUM_Z}, got {}", + self.num_z, + ); + + assert_eq!( + philen, PHILEN, + "LTIMES const-array preload version requires PHILEN = {PHILEN}, got {philen}", + ); + + assert_eq!( + elllen, ELLLEN, + "LTIMES const-array preload version requires ELLLEN = {ELLLEN}, got {elllen}", + ); + + assert_eq!( + psilen, PSILEN, + "LTIMES const-array preload version requires PSILEN = {PSILEN}, got {psilen}", + ); + unsafe { self.phidat = alloc_and_init_data_const(philen, to_real(0.0)); self.elldat = alloc_and_init_data(elllen); self.psidat = alloc_and_init_data(psilen); } + + let phidat_ref: &'a mut [Real; PHILEN] = + unsafe { &mut *(self.phidat as *mut [Real; PHILEN]) }; + + let elldat_ref: &'a [Real; ELLLEN] = unsafe { &*(self.elldat as *const [Real; ELLLEN]) }; + + let psidat_ref: &'a [Real; PSILEN] = unsafe { &*(self.psidat as *const [Real; PSILEN]) }; + + self.p_phidat = Some(preload_mut(phidat_ref)); + self.p_elldat = Some(preload(elldat_ref)); + self.p_psidat = Some(preload(psidat_ref)); } fn run_kernel(&mut self) { @@ -101,20 +148,41 @@ impl KernelBase for LTimes { let grid_y = NUM_G.div_ceil(g_block); let grid_z = num_z.div_ceil(z_block); - core::intrinsics::offload::<_, _, ()>( - _ltimes, - [grid_x as u32, grid_y as u32, grid_z as u32], - [m_block as u32, g_block as u32, z_block as u32], - ( - self.phidat as *mut [Real; 390400], - self.elldat as *const [Real; 1600], - self.psidat as *const [Real; 999424], + let p_phidat = self + .p_phidat + .as_ref() + .expect("LTIMES phidat was not preloaded"); + + let _p_elldat = self + .p_elldat + .as_ref() + .expect("LTIMES elldat was not preloaded"); + + let _p_psidat = self + .p_psidat + .as_ref() + .expect("LTIMES psidat was not preloaded"); + + let mut phidat_reg = Region::<'_, _, Stride3D<32, 8, 1, 25, 32>>::from(p_phidat); + + let elldat_ref: &[Real; ELLLEN] = unsafe { &*(self.elldat as *const [Real; ELLLEN]) }; + + let psidat_ref: &[Real; PSILEN] = unsafe { &*(self.psidat as *const [Real; PSILEN]) }; + + offload! { + kernel = ltimes, + grid_dim = [grid_x as u32, grid_y as u32, grid_z as u32], + block_dim = [m_block as u32, g_block as u32, z_block as u32], + args = ( + phidat_reg, + elldat_ref, + psidat_ref, NUM_D, NUM_M, NUM_G, num_z, ), - ); + }; } fn update_checksum(&self) -> f64 { @@ -122,63 +190,67 @@ impl KernelBase for LTimes { unsafe { calc_checksum(self.phidat as *const Real, philen) } } - fn tear_down(&mut self) { + fn tear_down(&mut self) -> f64 { + // Drop output mapping first, so device writes are copied back. + drop(self.p_phidat.take()); + + let checksum = self.update_checksum(); + + // Drop read-only mappings before freeing host allocations. + drop(self.p_elldat.take()); + drop(self.p_psidat.take()); + unsafe { - free(self.phidat); - free(self.elldat); - free(self.psidat); - self.phidat = core::ptr::null_mut(); - self.elldat = core::ptr::null_mut(); - self.psidat = core::ptr::null_mut(); + if !self.phidat.is_null() { + free(self.phidat); + self.phidat = core::ptr::null_mut(); + } + + if !self.elldat.is_null() { + free(self.elldat); + self.elldat = core::ptr::null_mut(); + } + + if !self.psidat.is_null() { + free(self.psidat); + self.psidat = core::ptr::null_mut(); + } } + self.num_z = 0; - } -} -#[cfg(target_os = "linux")] -unsafe extern "C" { - pub fn _ltimes( - phi: *mut [Real; 390400], - ell: &[Real; 1600], - psi: &[Real; 999424], - num_d: usize, - num_m: usize, - num_g: usize, - num_z: usize, - ); + checksum + } } #[cfg(not(target_os = "linux"))] use crate::common::types::Real; -#[cfg(not(target_os = "linux"))] -#[unsafe(no_mangle)] -#[rustc_offload_kernel] -pub unsafe extern "gpu-kernel" fn _ltimes( - phi: *mut [Real; 390400], - ell: &[Real; 1600], - psi: &[Real; 999424], +#[offload_kernel] +fn ltimes( + mut phi: Region>, + ell: &[Real; ELLLEN], + psi: &[Real; PSILEN], num_d: usize, num_m: usize, num_g: usize, num_z: usize, ) { - let num_m = NUM_M; - let num_g = NUM_G; - let num_d = NUM_D; - - let m = (block_idx_x() * 32 + thread_idx_x()) as usize; - let g = (block_idx_y() * 8 + thread_idx_y()) as usize; - let z = (block_idx_z() * 1 + thread_idx_z()) as usize; + let m = unsafe { (block_idx_x() * 32 + thread_idx_x()) as usize }; + let g = unsafe { (block_idx_y() * 8 + thread_idx_y()) as usize }; + let z = unsafe { (block_idx_z() * 1 + thread_idx_z()) as usize }; if m < num_m && g < num_g && z < num_z { - let phi_idx = m + num_m * (g + num_g * z); - - for d in 0..num_d { - let ell_idx = d + num_d * m; - let psi_idx = d + num_d * (g + num_g * z); + if let Some(v) = phi.get_mut() { + //#[rustc_unroll(8)] + for d in 0..num_d { + let ell_idx = d + num_d * m; + let psi_idx = d + num_d * (g + num_g * z); - (*phi)[phi_idx] += (*ell)[ell_idx] * (*psi)[psi_idx]; + unsafe { + *v += *ell.get_unchecked(ell_idx) * *psi.get_unchecked(psi_idx); + } + } } } } diff --git a/src/apps/mass3dea.rs b/src/apps/mass3dea.rs index 605c01d..caf5f8c 100644 --- a/src/apps/mass3dea.rs +++ b/src/apps/mass3dea.rs @@ -5,25 +5,16 @@ const DEFAULT_REPS: u32 = 1; const EA_MAT: usize = MEA_D1D * MEA_D1D * MEA_D1D * MEA_D1D * MEA_D1D * MEA_D1D; +#[cfg(target_arch = "amdgpu")] +use core::arch::amdgpu::{ + _syncthreads, workgroup_id_x as block_idx_x, workitem_id_x as thread_idx_x, + workitem_id_y as thread_idx_y, workitem_id_z as thread_idx_z, +}; #[cfg(target_arch = "nvptx64")] use core::arch::nvptx::{ _block_idx_x as block_idx_x, _syncthreads, _thread_idx_x as thread_idx_x, _thread_idx_y as thread_idx_y, _thread_idx_z as thread_idx_z, }; -#[cfg(target_arch = "amdgpu")] -#[allow(improper_ctypes)] -unsafe extern "C" { - #[link_name = "llvm.amdgcn.workgroup.id.x"] - fn block_idx_x() -> i32; - #[link_name = "llvm.amdgcn.workitem.id.x"] - fn thread_idx_x() -> i32; - #[link_name = "llvm.amdgcn.workitem.id.y"] - fn thread_idx_y() -> i32; - #[link_name = "llvm.amdgcn.workitem.id.z"] - fn thread_idx_z() -> i32; - #[link_name = "llvm.amdgcn.s.barrier"] - fn _syncthreads(); -} #[cfg(target_os = "linux")] use crate::common::data_utils::{alloc_and_init_data_const, calc_checksum, free}; diff --git a/src/apps/matvec_3d_stencil.rs b/src/apps/matvec_3d_stencil.rs index 35b826c..1565987 100644 --- a/src/apps/matvec_3d_stencil.rs +++ b/src/apps/matvec_3d_stencil.rs @@ -1,33 +1,44 @@ -pub const N_DEFAULT: usize = 1000000; +pub const N_DEFAULT: usize = 1_000_000; const DEFAULT_REPS: u32 = 100; +const M_DEFAULT: usize = 100; +const NPNL: usize = 2; +const NPNR: usize = 1; + +const DIM_SIZE: usize = M_DEFAULT + 1; +const NN_DIM: usize = DIM_SIZE + NPNL + NPNR; +const NNALLS: usize = NN_DIM * NN_DIM * NN_DIM; + +use core::offload::offload_kernel; + +#[cfg(target_arch = "amdgpu")] +use core::arch::amdgpu::{workgroup_id_x as block_idx_x, workitem_id_x as thread_idx_x}; + #[cfg(target_arch = "nvptx64")] use core::arch::nvptx::{_block_idx_x as block_idx_x, _thread_idx_x as thread_idx_x}; -#[cfg(target_arch = "amdgpu")] -#[allow(improper_ctypes)] -unsafe extern "C" { - #[link_name = "llvm.amdgcn.workgroup.id.x"] - fn block_idx_x() -> i32; - #[link_name = "llvm.amdgcn.workitem.id.x"] - fn thread_idx_x() -> i32; -} +#[cfg(target_os = "linux")] +use rustc_offload_frontend::offload; + +#[cfg(target_os = "linux")] +use core::offload::offload::{preload, preload_mut, Preload, PreloadMut}; #[cfg(target_os = "linux")] use crate::common::data_utils::{ - alloc_and_init_data, alloc_and_init_data_const, calc_checksum, calc_multiplier, free, - inc_data_init_count, + alloc_and_init_data, alloc_and_init_data_const, calc_checksum, free, inc_data_init_count, }; + #[cfg(target_os = "linux")] use crate::common::kernel_base::KernelBase; + #[cfg(target_os = "linux")] use crate::kernel_name; #[cfg(target_os = "linux")] -use crate::common::types::{Real, to_real}; +use crate::common::types::{to_real, Real}; #[cfg(target_os = "linux")] -pub struct Matvec3DStencil { +pub struct Matvec3DStencil<'a> { n: usize, x: *mut Real, b: *mut Real, @@ -37,10 +48,15 @@ pub struct Matvec3DStencil { kp: usize, fpz: usize, lpz: usize, + + p_x: Option>, + p_b: Option>, + p_matrix: [Option>; 14], + p_real_zones: Option>, } #[cfg(target_os = "linux")] -impl Matvec3DStencil { +impl<'a> Matvec3DStencil<'a> { pub const INIT: Self = Matvec3DStencil { n: 0, x: core::ptr::null_mut(), @@ -51,11 +67,16 @@ impl Matvec3DStencil { kp: 0, fpz: 0, lpz: 0, + + p_x: None, + p_b: None, + p_matrix: [const { None }; 14], + p_real_zones: None, }; } #[cfg(target_os = "linux")] -impl KernelBase for Matvec3DStencil { +impl<'a> KernelBase for Matvec3DStencil<'a> { fn name(&self) -> &'static str { kernel_name!("MATVEC_3D_STENCIL") } @@ -69,15 +90,29 @@ impl KernelBase for Matvec3DStencil { } fn setup(&mut self) { - self.n = N_DEFAULT; - let m: usize = 100; - let npnl = 2; - let npnr = 1; + self.n = self.default_problem_size(); + + assert_eq!( + self.n, N_DEFAULT, + "MATVEC_3D_STENCIL const-array preload version requires n = {N_DEFAULT}, got {}", + self.n + ); + + let m: usize = M_DEFAULT; + let npnl = NPNL; + let npnr = NPNR; + let dim_size = m + 1; let nn_dim = dim_size + npnl + npnr; + let nnalls = nn_dim * nn_dim * nn_dim; + + assert_eq!( + nnalls, NNALLS, + "MATVEC_3D_STENCIL const-array preload version requires NNALLS = {NNALLS}, got {nnalls}", + ); + self.jp = nn_dim; self.kp = nn_dim * nn_dim; - let nnalls = nn_dim * nn_dim * nn_dim; let imin = npnl; let imax = npnl + dim_size - 1; @@ -93,17 +128,19 @@ impl KernelBase for Matvec3DStencil { self.lpz = (kmax - 1 + npzr) * self.kp + (jmax - 1 + npzr) * self.jp + (imax - 1 + npzr); unsafe { - self.b = alloc_and_init_data_const(nnalls, to_real(0.0)); - self.x = alloc_and_init_data(nnalls); + self.b = alloc_and_init_data_const(NNALLS, to_real(0.0)); + self.x = alloc_and_init_data(NNALLS); for i in 0..14 { - self.matrix[i] = alloc_and_init_data(nnalls); + self.matrix[i] = alloc_and_init_data(NNALLS); } self.real_zones = crate::common::data_utils::alloc::(self.n); + for i in 0..self.n { *self.real_zones.add(i) = u64::MAX; } + inc_data_init_count(); let j_stride = imax - imin; @@ -114,11 +151,30 @@ impl KernelBase for Matvec3DStencil { for i in imin..imax { let iz = i + j * self.jp + k * self.kp; let il = (i - imin) + (j - jmin) * j_stride + (k - kmin) * k_stride; + *self.real_zones.add(il) = iz as u64; } } } } + + let x_ref: &'a [Real; NNALLS] = unsafe { &*(self.x as *const [Real; NNALLS]) }; + + let b_ref: &'a mut [Real; NNALLS] = unsafe { &mut *(self.b as *mut [Real; NNALLS]) }; + + let real_zones_ref: &'a [u64; N_DEFAULT] = + unsafe { &*(self.real_zones as *const [u64; N_DEFAULT]) }; + + self.p_x = Some(preload(x_ref)); + self.p_b = Some(preload_mut(b_ref)); + self.p_real_zones = Some(preload(real_zones_ref)); + + for i in 0..14 { + let matrix_ref: &'a [Real; NNALLS] = + unsafe { &*(self.matrix[i] as *const [Real; NNALLS]) }; + + self.p_matrix[i] = Some(preload(matrix_ref)); + } } fn run_kernel(&mut self) { @@ -126,111 +182,130 @@ impl KernelBase for Matvec3DStencil { let jp = self.jp; let kp = self.kp; - core::intrinsics::offload::<_, _, ()>( - _matvec3dstencil, - [n.div_ceil(256) as u32, 1, 1], - [256, 1, 1], - ( - self.x as *const [Real; 1124864], - self.b as *mut [Real; 1124864], - self.matrix[0] as *const [Real; 1124864], - self.matrix[1] as *const [Real; 1124864], - self.matrix[2] as *const [Real; 1124864], - self.matrix[3] as *const [Real; 1124864], - self.matrix[4] as *const [Real; 1124864], - self.matrix[5] as *const [Real; 1124864], - self.matrix[6] as *const [Real; 1124864], - self.matrix[7] as *const [Real; 1124864], - self.matrix[8] as *const [Real; 1124864], - self.matrix[9] as *const [Real; 1124864], - self.matrix[10] as *const [Real; 1124864], - self.matrix[11] as *const [Real; 1124864], - self.matrix[12] as *const [Real; 1124864], - self.matrix[13] as *const [Real; 1124864], - self.real_zones as *const [u64; 1000000], + self.p_x + .as_ref() + .expect("MATVEC_3D_STENCIL x was not preloaded"); + + self.p_b + .as_ref() + .expect("MATVEC_3D_STENCIL b was not preloaded"); + + self.p_real_zones + .as_ref() + .expect("MATVEC_3D_STENCIL real_zones was not preloaded"); + + for p in &self.p_matrix { + p.as_ref() + .expect("MATVEC_3D_STENCIL matrix entry was not preloaded"); + } + + offload! { + kernel = matvec3dstencil, + grid_dim = [n.div_ceil(256) as u32, 1, 1], + block_dim = [256, 1, 1], + args = ( + self.x as *const [Real; NNALLS], + self.b as *mut [Real; NNALLS], + self.matrix[0] as *const [Real; NNALLS], + self.matrix[1] as *const [Real; NNALLS], + self.matrix[2] as *const [Real; NNALLS], + self.matrix[3] as *const [Real; NNALLS], + self.matrix[4] as *const [Real; NNALLS], + self.matrix[5] as *const [Real; NNALLS], + self.matrix[6] as *const [Real; NNALLS], + self.matrix[7] as *const [Real; NNALLS], + self.matrix[8] as *const [Real; NNALLS], + self.matrix[9] as *const [Real; NNALLS], + self.matrix[10] as *const [Real; NNALLS], + self.matrix[11] as *const [Real; NNALLS], + self.matrix[12] as *const [Real; NNALLS], + self.matrix[13] as *const [Real; NNALLS], + self.real_zones as *const [u64; N_DEFAULT], jp, kp, n, ), - ); + }; } fn update_checksum(&self) -> f64 { - unsafe { calc_checksum(self.b as *const Real, 1124864) } + unsafe { calc_checksum(self.b as *const Real, NNALLS) } } - fn tear_down(&mut self) { + fn tear_down(&mut self) -> f64 { + // Drop preloads first. Dropping p_b copies device output back to host. + drop(self.p_b.take()); + drop(self.p_x.take()); + drop(self.p_real_zones.take()); + + for p in &mut self.p_matrix { + drop(p.take()); + } + + let checksum = self.update_checksum(); + unsafe { - free(self.x); - free(self.b); + if !self.x.is_null() { + free(self.x); + self.x = core::ptr::null_mut(); + } + + if !self.b.is_null() { + free(self.b); + self.b = core::ptr::null_mut(); + } + for i in 0..14 { - free(self.matrix[i]); - self.matrix[i] = core::ptr::null_mut(); + if !self.matrix[i].is_null() { + free(self.matrix[i]); + self.matrix[i] = core::ptr::null_mut(); + } + } + + if !self.real_zones.is_null() { + free(self.real_zones); + self.real_zones = core::ptr::null_mut(); } - free(self.real_zones); - self.x = core::ptr::null_mut(); - self.b = core::ptr::null_mut(); - self.real_zones = core::ptr::null_mut(); } + self.n = 0; - } -} + self.jp = 0; + self.kp = 0; + self.fpz = 0; + self.lpz = 0; -#[cfg(target_os = "linux")] -unsafe extern "C" { - pub fn _matvec3dstencil( - x: *const [Real; 1124864], - b: *mut [Real; 1124864], - m0: *const [Real; 1124864], - m1: *const [Real; 1124864], - m2: *const [Real; 1124864], - m3: *const [Real; 1124864], - m4: *const [Real; 1124864], - m5: *const [Real; 1124864], - m6: *const [Real; 1124864], - m7: *const [Real; 1124864], - m8: *const [Real; 1124864], - m9: *const [Real; 1124864], - m10: *const [Real; 1124864], - m11: *const [Real; 1124864], - m12: *const [Real; 1124864], - m13: *const [Real; 1124864], - real_zones: *const [u64; 1000000], - jp: usize, - kp: usize, - n: usize, - ); + checksum + } } #[cfg(not(target_os = "linux"))] use crate::common::types::Real; -#[cfg(not(target_os = "linux"))] -#[unsafe(no_mangle)] -#[rustc_offload_kernel] -pub unsafe extern "gpu-kernel" fn _matvec3dstencil( - x: *const [Real; 1124864], - b: *mut [Real; 1124864], - m0: *const [Real; 1124864], - m1: *const [Real; 1124864], - m2: *const [Real; 1124864], - m3: *const [Real; 1124864], - m4: *const [Real; 1124864], - m5: *const [Real; 1124864], - m6: *const [Real; 1124864], - m7: *const [Real; 1124864], - m8: *const [Real; 1124864], - m9: *const [Real; 1124864], - m10: *const [Real; 1124864], - m11: *const [Real; 1124864], - m12: *const [Real; 1124864], - m13: *const [Real; 1124864], - real_zones: *const [u64; 1000000], +#[offload_kernel] +fn matvec3dstencil( + x: *const [Real; NNALLS], + b: *mut [Real; NNALLS], + m0: *const [Real; NNALLS], + m1: *const [Real; NNALLS], + m2: *const [Real; NNALLS], + m3: *const [Real; NNALLS], + m4: *const [Real; NNALLS], + m5: *const [Real; NNALLS], + m6: *const [Real; NNALLS], + m7: *const [Real; NNALLS], + m8: *const [Real; NNALLS], + m9: *const [Real; NNALLS], + m10: *const [Real; NNALLS], + m11: *const [Real; NNALLS], + m12: *const [Real; NNALLS], + m13: *const [Real; NNALLS], + real_zones: *const [u64; N_DEFAULT], jp: usize, kp: usize, n: usize, ) { let i = unsafe { (block_idx_x() * 256 + thread_idx_x()) as usize }; + if i >= n { return; } @@ -241,6 +316,7 @@ pub unsafe extern "gpu-kernel" fn _matvec3dstencil( unsafe { let mut b0 = (*m0)[iz] * (*x)[(iz as isize - 1 - s_jp - s_kp) as usize]; + b0 += (*m1)[iz] * (*x)[(iz as isize - s_jp - s_kp) as usize]; b0 += (*m2)[iz] * (*x)[(iz as isize + 1 - s_jp - s_kp) as usize]; b0 += (*m3)[iz] * (*x)[(iz as isize - 1 - s_kp) as usize]; diff --git a/src/apps/mod.rs b/src/apps/mod.rs index f8bb4fe..ea37ba8 100644 --- a/src/apps/mod.rs +++ b/src/apps/mod.rs @@ -1,3 +1,4 @@ +#![feature(stmt_expr_attributes)] #[cfg(feature = "del_dot_vec_2d")] pub mod del_dot_vec_2d; #[cfg(feature = "energy")] diff --git a/src/apps/pressure.rs b/src/apps/pressure.rs index 582f039..83fe205 100644 --- a/src/apps/pressure.rs +++ b/src/apps/pressure.rs @@ -1,32 +1,31 @@ -pub const N_DEFAULT: usize = 1000000; +pub const N_DEFAULT: usize = 1_000_000; const DEFAULT_REPS: u32 = 700; -#[cfg(target_arch = "nvptx64")] -use core::arch::nvptx::{_block_idx_x as block_idx_x, _thread_idx_x as thread_idx_x}; +use core::offload::offload_kernel; +use rustc_offload_frontend::partition::{PartitioningStrategy, Region, Stride1D}; -#[cfg(target_arch = "amdgpu")] -#[allow(improper_ctypes)] -unsafe extern "C" { - #[link_name = "llvm.amdgcn.workgroup.id.x"] - fn block_idx_x() -> i32; - #[link_name = "llvm.amdgcn.workitem.id.x"] - fn thread_idx_x() -> i32; -} +#[cfg(target_os = "linux")] +use rustc_offload_frontend::offload; + +#[cfg(target_os = "linux")] +use core::offload::offload::{preload, preload_mut, Preload, PreloadMut}; #[cfg(target_os = "linux")] use crate::common::data_utils::{ alloc_and_init_data, alloc_and_init_data_const, calc_checksum, free, init_data_scalar, }; + #[cfg(target_os = "linux")] use crate::common::kernel_base::KernelBase; + #[cfg(target_os = "linux")] use crate::kernel_name; #[cfg(target_os = "linux")] -use crate::common::types::{Real, to_real}; +use crate::common::types::{to_real, Real}; #[cfg(target_os = "linux")] -pub struct Pressure { +pub struct Pressure<'a> { n: usize, compression: *mut Real, bvc: *mut Real, @@ -37,10 +36,16 @@ pub struct Pressure { p_cut: Real, pmin: Real, eosvmax: Real, + + p_compression: Option>, + p_bvc: Option>, + p_p_new: Option>, + p_e_old: Option>, + p_vnewc: Option>, } #[cfg(target_os = "linux")] -impl Pressure { +impl<'a> Pressure<'a> { pub const INIT: Self = Pressure { n: 0, compression: core::ptr::null_mut(), @@ -52,11 +57,17 @@ impl Pressure { p_cut: to_real(1.0e-7), pmin: to_real(1.0e-10), eosvmax: to_real(1.0e+10), + + p_compression: None, + p_bvc: None, + p_p_new: None, + p_e_old: None, + p_vnewc: None, }; } #[cfg(target_os = "linux")] -impl KernelBase for Pressure { +impl<'a> KernelBase for Pressure<'a> { fn name(&self) -> &'static str { kernel_name!("PRESSURE") } @@ -70,7 +81,13 @@ impl KernelBase for Pressure { } fn setup(&mut self) { - self.n = N_DEFAULT; + self.n = self.default_problem_size(); + + assert_eq!( + self.n, N_DEFAULT, + "PRESSURE const-array preload version requires n = {N_DEFAULT}, got {}", + self.n + ); unsafe { self.compression = alloc_and_init_data(self.n); @@ -84,6 +101,27 @@ impl KernelBase for Pressure { self.pmin = init_data_scalar(); self.eosvmax = init_data_scalar(); } + + let compression_ref: &'a [Real; N_DEFAULT] = + unsafe { &*(self.compression as *const [Real; N_DEFAULT]) }; + + let bvc_ref: &'a mut [Real; N_DEFAULT] = + unsafe { &mut *(self.bvc as *mut [Real; N_DEFAULT]) }; + + let p_new_ref: &'a mut [Real; N_DEFAULT] = + unsafe { &mut *(self.p_new as *mut [Real; N_DEFAULT]) }; + + let e_old_ref: &'a [Real; N_DEFAULT] = + unsafe { &*(self.e_old as *const [Real; N_DEFAULT]) }; + + let vnewc_ref: &'a [Real; N_DEFAULT] = + unsafe { &*(self.vnewc as *const [Real; N_DEFAULT]) }; + + self.p_compression = Some(preload(compression_ref)); + self.p_bvc = Some(preload_mut(bvc_ref)); + self.p_p_new = Some(preload_mut(p_new_ref)); + self.p_e_old = Some(preload(e_old_ref)); + self.p_vnewc = Some(preload(vnewc_ref)); } fn run_kernel(&mut self) { @@ -91,24 +129,46 @@ impl KernelBase for Pressure { let grid = [n.div_ceil(256) as u32, 1, 1]; let block = [256, 1, 1]; - core::intrinsics::offload::<_, _, ()>( - _pressure_calc1, - grid, - block, - ( - self.bvc as *mut [Real; N_DEFAULT], + self.p_compression + .as_ref() + .expect("PRESSURE compression was not preloaded"); + + let p_bvc = self.p_bvc.as_ref().expect("PRESSURE bvc was not preloaded"); + + let p_p_new = self + .p_p_new + .as_ref() + .expect("PRESSURE p_new was not preloaded"); + + self.p_e_old + .as_ref() + .expect("PRESSURE e_old was not preloaded"); + + self.p_vnewc + .as_ref() + .expect("PRESSURE vnewc was not preloaded"); + + let mut bvc_reg = Region::<'_, _, Stride1D<256>>::from(p_bvc); + let mut p_new_reg = Region::<'_, _, Stride1D<256>>::from(p_p_new); + + offload! { + kernel = pressure_calc1, + grid_dim = grid, + block_dim = block, + args = ( + bvc_reg, self.compression as *const [Real; N_DEFAULT], self.cls, n, ), - ); + }; - core::intrinsics::offload::<_, _, ()>( - _pressure_calc2, - grid, - block, - ( - self.p_new as *mut [Real; N_DEFAULT], + offload! { + kernel = pressure_calc2, + grid_dim = grid, + block_dim = block, + args = ( + p_new_reg, self.bvc as *const [Real; N_DEFAULT], self.e_old as *const [Real; N_DEFAULT], self.vnewc as *const [Real; N_DEFAULT], @@ -117,76 +177,79 @@ impl KernelBase for Pressure { self.pmin, n, ), - ); + }; } fn update_checksum(&self) -> f64 { unsafe { calc_checksum(self.p_new as *const Real, self.n) } } - fn tear_down(&mut self) { + fn tear_down(&mut self) -> f64 { + // Drop preloads first. Dropping p_bvc / p_p_new makes device writes visible on host. + drop(self.p_bvc.take()); + drop(self.p_p_new.take()); + + drop(self.p_compression.take()); + drop(self.p_e_old.take()); + drop(self.p_vnewc.take()); + + let checksum = self.update_checksum(); + unsafe { - free(self.compression); - free(self.bvc); - free(self.p_new); - free(self.e_old); - free(self.vnewc); - self.compression = core::ptr::null_mut(); - self.bvc = core::ptr::null_mut(); - self.p_new = core::ptr::null_mut(); - self.e_old = core::ptr::null_mut(); - self.vnewc = core::ptr::null_mut(); + if !self.compression.is_null() { + free(self.compression); + self.compression = core::ptr::null_mut(); + } + + if !self.bvc.is_null() { + free(self.bvc); + self.bvc = core::ptr::null_mut(); + } + + if !self.p_new.is_null() { + free(self.p_new); + self.p_new = core::ptr::null_mut(); + } + + if !self.e_old.is_null() { + free(self.e_old); + self.e_old = core::ptr::null_mut(); + } + + if !self.vnewc.is_null() { + free(self.vnewc); + self.vnewc = core::ptr::null_mut(); + } } + self.n = 0; - } -} -#[cfg(target_os = "linux")] -unsafe extern "C" { - pub fn _pressure_calc1( - bvc: *mut [Real; N_DEFAULT], - compression: *const [Real; N_DEFAULT], - cls: Real, - n: usize, - ); - - pub fn _pressure_calc2( - p_new: *mut [Real; N_DEFAULT], - bvc: *const [Real; N_DEFAULT], - e_old: *const [Real; N_DEFAULT], - vnewc: *const [Real; N_DEFAULT], - p_cut: Real, - eosvmax: Real, - pmin: Real, - n: usize, - ); + checksum + } } #[cfg(not(target_os = "linux"))] use crate::common::types::Real; -#[cfg(not(target_os = "linux"))] -#[unsafe(no_mangle)] -#[rustc_offload_kernel] -pub unsafe extern "gpu-kernel" fn _pressure_calc1( - bvc: *mut [Real; N_DEFAULT], +#[offload_kernel] +fn pressure_calc1( + mut bvc: Region>, compression: *const [Real; N_DEFAULT], cls: Real, n: usize, ) { - let i = unsafe { (block_idx_x() * 256 + thread_idx_x()) as usize }; - if i < n { + let i = Stride1D::<256>::index(); + + if let Some(v) = bvc.get_mut() { unsafe { - (*bvc)[i] = cls * ((*compression)[i] + Real::from(1.0)); + *v = cls * ((*compression)[i] + Real::from(1.0)); } } } -#[cfg(not(target_os = "linux"))] -#[unsafe(no_mangle)] -#[rustc_offload_kernel] -pub unsafe extern "gpu-kernel" fn _pressure_calc2( - p_new: *mut [Real; N_DEFAULT], +#[offload_kernel] +fn pressure_calc2( + mut p_new: Region>, bvc: *const [Real; N_DEFAULT], e_old: *const [Real; N_DEFAULT], vnewc: *const [Real; N_DEFAULT], @@ -195,22 +258,25 @@ pub unsafe extern "gpu-kernel" fn _pressure_calc2( pmin: Real, n: usize, ) { - let i = unsafe { (block_idx_x() * 256 + thread_idx_x()) as usize }; - if i < n { + let i = Stride1D::<256>::index(); + + if let Some(v) = p_new.get_mut() { unsafe { let mut p = (*bvc)[i] * (*e_old)[i]; if p.abs() < p_cut { p = Real::from(0.0); } + if (*vnewc)[i] >= eosvmax { p = Real::from(0.0); } + if p < pmin { p = pmin; } - (*p_new)[i] = p; + *v = p; } } } diff --git a/src/apps/vol3d.rs b/src/apps/vol3d.rs index 7cf5877..dca4d74 100644 --- a/src/apps/vol3d.rs +++ b/src/apps/vol3d.rs @@ -1,30 +1,37 @@ -pub const N_DEFAULT: usize = 1000000; +pub const N_DEFAULT: usize = 1_000_000; const DEFAULT_REPS: u32 = 100; -#[cfg(target_arch = "nvptx64")] -use core::arch::nvptx::{_block_idx_x as block_idx_x, _thread_idx_x as thread_idx_x}; +const M_DEFAULT: usize = 100; +const NPNL: usize = 2; +const NPNR: usize = 1; -#[cfg(target_arch = "amdgpu")] -#[allow(improper_ctypes)] -unsafe extern "C" { - #[link_name = "llvm.amdgcn.workgroup.id.x"] - fn block_idx_x() -> i32; - #[link_name = "llvm.amdgcn.workitem.id.x"] - fn thread_idx_x() -> i32; -} +const DIM_SIZE: usize = M_DEFAULT + 1; +const NN_DIM: usize = DIM_SIZE + NPNL + NPNR; +const NNALLS: usize = NN_DIM * NN_DIM * NN_DIM; + +use core::offload::offload_kernel; +use rustc_offload_frontend::partition::{OffsetStride1D, PartitioningStrategy, Region}; + +#[cfg(target_os = "linux")] +use rustc_offload_frontend::offload; + +#[cfg(target_os = "linux")] +use core::offload::offload::{preload, preload_mut, Preload, PreloadMut}; #[cfg(target_os = "linux")] use crate::common::data_utils::{alloc_and_init_data_const, calc_checksum, free}; + #[cfg(target_os = "linux")] use crate::common::kernel_base::KernelBase; + #[cfg(target_os = "linux")] use crate::kernel_name; #[cfg(target_os = "linux")] -use crate::common::types::{Real, to_real}; +use crate::common::types::{to_real, Real}; #[cfg(target_os = "linux")] -pub struct Vol3D { +pub struct Vol3D<'a> { n: usize, x: *mut Real, y: *mut Real, @@ -35,10 +42,15 @@ pub struct Vol3D { kp: usize, fpz: usize, lpz: usize, + + p_x: Option>, + p_y: Option>, + p_z: Option>, + p_vol: Option>, } #[cfg(target_os = "linux")] -impl Vol3D { +impl<'a> Vol3D<'a> { pub const INIT: Self = Vol3D { n: 0, x: core::ptr::null_mut(), @@ -50,11 +62,16 @@ impl Vol3D { kp: 0, fpz: 0, lpz: 0, + + p_x: None, + p_y: None, + p_z: None, + p_vol: None, }; } #[cfg(target_os = "linux")] -impl KernelBase for Vol3D { +impl<'a> KernelBase for Vol3D<'a> { fn name(&self) -> &'static str { kernel_name!("VOL3D") } @@ -68,15 +85,29 @@ impl KernelBase for Vol3D { } fn setup(&mut self) { - self.n = N_DEFAULT; - let m: usize = 100; - let npnl = 2; - let npnr = 1; + self.n = self.default_problem_size(); + + assert_eq!( + self.n, N_DEFAULT, + "VOL3D const-array preload version requires n = {N_DEFAULT}, got {}", + self.n + ); + + let m: usize = M_DEFAULT; + let npnl = NPNL; + let npnr = NPNR; + let dim_size = m + 1; let nn_dim = dim_size + npnl + npnr; + let nnalls = nn_dim * nn_dim * nn_dim; + + assert_eq!( + nnalls, NNALLS, + "VOL3D const-array preload version requires NNALLS = {NNALLS}, got {nnalls}", + ); + self.jp = nn_dim; self.kp = nn_dim * nn_dim; - let nnalls = nn_dim * nn_dim * nn_dim; let imin = npnl; let imax = npnl + dim_size - 1; @@ -92,10 +123,10 @@ impl KernelBase for Vol3D { self.lpz = (kmax - 1 + npzr) * self.kp + (jmax - 1 + npzr) * self.jp + (imax - 1 + npzr); unsafe { - self.x = alloc_and_init_data_const(nnalls, to_real(0.0)); - self.y = alloc_and_init_data_const(nnalls, to_real(0.0)); - self.z = alloc_and_init_data_const(nnalls, to_real(0.0)); - self.vol = alloc_and_init_data_const(nnalls, to_real(0.0)); + self.x = alloc_and_init_data_const(NNALLS, to_real(0.0)); + self.y = alloc_and_init_data_const(NNALLS, to_real(0.0)); + self.z = alloc_and_init_data_const(NNALLS, to_real(0.0)); + self.vol = alloc_and_init_data_const(NNALLS, to_real(0.0)); let dx = 0.3; let dy = 0.2; @@ -105,6 +136,7 @@ impl KernelBase for Vol3D { for j in (jmin - npnl)..(jmax + 1 + npnr) { for i in (imin - npnl)..(imax + 1 + npnr) { let inn = i + j * self.jp + k * self.kp; + *self.x.add(inn) = to_real(i as f64 * dx); *self.y.add(inn) = to_real(j as f64 * dy); *self.z.add(inn) = to_real(k as f64 * dz); @@ -112,6 +144,19 @@ impl KernelBase for Vol3D { } } } + + let x_ref: &'a [Real; NNALLS] = unsafe { &*(self.x as *const [Real; NNALLS]) }; + + let y_ref: &'a [Real; NNALLS] = unsafe { &*(self.y as *const [Real; NNALLS]) }; + + let z_ref: &'a [Real; NNALLS] = unsafe { &*(self.z as *const [Real; NNALLS]) }; + + let vol_ref: &'a mut [Real; NNALLS] = unsafe { &mut *(self.vol as *mut [Real; NNALLS]) }; + + self.p_x = Some(preload(x_ref)); + self.p_y = Some(preload(y_ref)); + self.p_z = Some(preload(z_ref)); + self.p_vol = Some(preload_mut(vol_ref)); } fn run_kernel(&mut self) { @@ -121,139 +166,161 @@ impl KernelBase for Vol3D { let kp = self.kp; let count = lpz + 1 - fpz; - core::intrinsics::offload::<_, _, ()>( - _vol3d, - [((count + 255) / 256) as u32, 1, 1], - [256, 1, 1], - ( - self.x as *const [Real; 1124864], - self.y as *const [Real; 1124864], - self.z as *const [Real; 1124864], - self.vol as *mut [Real; 1124864], + self.p_x.as_ref().expect("VOL3D x was not preloaded"); + + self.p_y.as_ref().expect("VOL3D y was not preloaded"); + + self.p_z.as_ref().expect("VOL3D z was not preloaded"); + + let p_vol = self.p_vol.as_ref().expect("VOL3D vol was not preloaded"); + + let mut vol_reg = Region::<'_, _, OffsetStride1D<256>>::from(p_vol); + + offload! { + kernel = vol3d, + grid_dim = [count.div_ceil(256) as u32, 1, 1], + block_dim = [256, 1, 1], + args = ( + self.x as *const [Real; NNALLS], + self.y as *const [Real; NNALLS], + self.z as *const [Real; NNALLS], + vol_reg, self.vnormq, jp, kp, fpz, lpz, ), - ); + }; } fn update_checksum(&self) -> f64 { - unsafe { calc_checksum(self.vol as *const Real, 1124864) } + unsafe { calc_checksum(self.vol as *const Real, NNALLS) } } - fn tear_down(&mut self) { + fn tear_down(&mut self) -> f64 { + // Drop preloads first. Dropping p_vol copies device output back to host. + drop(self.p_vol.take()); + drop(self.p_x.take()); + drop(self.p_y.take()); + drop(self.p_z.take()); + + let checksum = self.update_checksum(); + unsafe { - free(self.x); - free(self.y); - free(self.z); - free(self.vol); - self.x = core::ptr::null_mut(); - self.y = core::ptr::null_mut(); - self.z = core::ptr::null_mut(); - self.vol = core::ptr::null_mut(); + if !self.x.is_null() { + free(self.x); + self.x = core::ptr::null_mut(); + } + + if !self.y.is_null() { + free(self.y); + self.y = core::ptr::null_mut(); + } + + if !self.z.is_null() { + free(self.z); + self.z = core::ptr::null_mut(); + } + + if !self.vol.is_null() { + free(self.vol); + self.vol = core::ptr::null_mut(); + } } + self.n = 0; - } -} + self.jp = 0; + self.kp = 0; + self.fpz = 0; + self.lpz = 0; -#[cfg(target_os = "linux")] -unsafe extern "C" { - pub fn _vol3d( - x: *const [Real; 1124864], - y: *const [Real; 1124864], - z: *const [Real; 1124864], - vol: *mut [Real; 1124864], - vnormq: Real, - jp: usize, - kp: usize, - fpz: usize, - lpz: usize, - ); + checksum + } } #[cfg(not(target_os = "linux"))] use crate::common::types::Real; -#[cfg(not(target_os = "linux"))] -#[unsafe(no_mangle)] -#[rustc_offload_kernel] -pub unsafe extern "gpu-kernel" fn _vol3d( - x: *const [Real; 1124864], - y: *const [Real; 1124864], - z: *const [Real; 1124864], - vol: *mut [Real; 1124864], +#[offload_kernel] +fn vol3d( + x: *const [Real; NNALLS], + y: *const [Real; NNALLS], + z: *const [Real; NNALLS], + mut vol: Region>, vnormq: Real, jp: usize, kp: usize, fpz: usize, lpz: usize, ) { - let idx = unsafe { (block_idx_x() * 256 + thread_idx_x()) as usize }; + let idx = OffsetStride1D::<256>::index(); let i = fpz + idx; + if i > lpz { return; } - unsafe { - let i0 = i; - let i1 = i + 1; - let i2 = i + jp; - let i3 = i + 1 + jp; - let i4 = i + kp; - let i5 = i + 1 + kp; - let i6 = i + jp + kp; - let i7 = i + 1 + jp + kp; - - let x71 = (*x)[i7] - (*x)[i1]; - let x72 = (*x)[i7] - (*x)[i2]; - let x74 = (*x)[i7] - (*x)[i4]; - let x30 = (*x)[i3] - (*x)[i0]; - let x50 = (*x)[i5] - (*x)[i0]; - let x60 = (*x)[i6] - (*x)[i0]; - - let y71 = (*y)[i7] - (*y)[i1]; - let y72 = (*y)[i7] - (*y)[i2]; - let y74 = (*y)[i7] - (*y)[i4]; - let y30 = (*y)[i3] - (*y)[i0]; - let y50 = (*y)[i5] - (*y)[i0]; - let y60 = (*y)[i6] - (*y)[i0]; - - let z71 = (*z)[i7] - (*z)[i1]; - let z72 = (*z)[i7] - (*z)[i2]; - let z74 = (*z)[i7] - (*z)[i4]; - let z30 = (*z)[i3] - (*z)[i0]; - let z50 = (*z)[i5] - (*z)[i0]; - let z60 = (*z)[i6] - (*z)[i0]; - - let mut xps = x71 + x60; - let mut yps = y71 + y60; - let mut zps = z71 + z60; - - let mut cyz = y72 * z30 - z72 * y30; - let mut czx = z72 * x30 - x72 * z30; - let mut cxy = x72 * y30 - y72 * x30; - let mut v = xps * cyz + yps * czx + zps * cxy; - - xps = x72 + x50; - yps = y72 + y50; - zps = z72 + z50; - - cyz = y74 * z60 - z74 * y60; - czx = z74 * x60 - x74 * z60; - cxy = x74 * y60 - y74 * x60; - v += xps * cyz + yps * czx + zps * cxy; - - xps = x74 + x30; - yps = y74 + y30; - zps = z74 + z30; - - cyz = y71 * z50 - z71 * y50; - czx = z71 * x50 - x71 * z50; - cxy = x71 * y50 - y71 * x50; - v += xps * cyz + yps * czx + zps * cxy; - - (*vol)[i] = v * vnormq; + if let Some(mut vvol) = vol.get_mut() { + unsafe { + let i0 = i; + let i1 = i + 1; + let i2 = i + jp; + let i3 = i + 1 + jp; + let i4 = i + kp; + let i5 = i + 1 + kp; + let i6 = i + jp + kp; + let i7 = i + 1 + jp + kp; + + let x71 = (*x)[i7] - (*x)[i1]; + let x72 = (*x)[i7] - (*x)[i2]; + let x74 = (*x)[i7] - (*x)[i4]; + let x30 = (*x)[i3] - (*x)[i0]; + let x50 = (*x)[i5] - (*x)[i0]; + let x60 = (*x)[i6] - (*x)[i0]; + + let y71 = (*y)[i7] - (*y)[i1]; + let y72 = (*y)[i7] - (*y)[i2]; + let y74 = (*y)[i7] - (*y)[i4]; + let y30 = (*y)[i3] - (*y)[i0]; + let y50 = (*y)[i5] - (*y)[i0]; + let y60 = (*y)[i6] - (*y)[i0]; + + let z71 = (*z)[i7] - (*z)[i1]; + let z72 = (*z)[i7] - (*z)[i2]; + let z74 = (*z)[i7] - (*z)[i4]; + let z30 = (*z)[i3] - (*z)[i0]; + let z50 = (*z)[i5] - (*z)[i0]; + let z60 = (*z)[i6] - (*z)[i0]; + + let mut xps = x71 + x60; + let mut yps = y71 + y60; + let mut zps = z71 + z60; + + let mut cyz = y72 * z30 - z72 * y30; + let mut czx = z72 * x30 - x72 * z30; + let mut cxy = x72 * y30 - y72 * x30; + let mut v = xps * cyz + yps * czx + zps * cxy; + + xps = x72 + x50; + yps = y72 + y50; + zps = z72 + z50; + + cyz = y74 * z60 - z74 * y60; + czx = z74 * x60 - x74 * z60; + cxy = x74 * y60 - y74 * x60; + v += xps * cyz + yps * czx + zps * cxy; + + xps = x74 + x30; + yps = y74 + y30; + zps = z74 + z30; + + cyz = y71 * z50 - z71 * y50; + czx = z71 * x50 - x71 * z50; + cxy = x71 * y50 - y71 * x50; + v += xps * cyz + yps * czx + zps * cxy; + + vvol.set(fpz, v * vnormq); + } } } diff --git a/src/common/executor.rs b/src/common/executor.rs index 23a41b6..d72a7d7 100644 --- a/src/common/executor.rs +++ b/src/common/executor.rs @@ -1,6 +1,7 @@ #[cfg(target_os = "linux")] extern crate libc; - +#[cfg(target_os = "linux")] +use std::offload::offload::offload_sync; #[cfg(target_os = "linux")] use crate::common::data_utils::reset_data_init_count; #[cfg(target_os = "linux")] @@ -14,6 +15,7 @@ pub struct KernelResult { pub problem_size: usize, pub reps: u32, pub total_s: f64, + pub launch_s: f64, pub checksum: f64, } @@ -36,27 +38,48 @@ impl<'a> Executor<'a> { out: &'r mut [core::mem::MaybeUninit; MAX_KERNELS], ) -> &'r [KernelResult] { let n = self.kernels.len().min(MAX_KERNELS); + + // warmup + //for (i, kernel) in self.kernels.iter_mut().enumerate().take(n) { + // let reps = kernel.default_reps(); + // kernel.setup(); + // for _ in 0..reps { + // kernel.run_kernel(); + // } + // let checksum = kernel.tear_down(); + //} for (i, kernel) in self.kernels.iter_mut().enumerate().take(n) { let reps = kernel.default_reps(); reset_data_init_count(); kernel.setup(); + //warmup + for _ in 0..50 { + kernel.run_kernel(); + //offload_sync(); + } + offload_sync(); let t0 = now_ns(); for _ in 0..reps { kernel.run_kernel(); + offload_sync(); } let t1 = now_ns(); + offload_sync(); + let t2 = now_ns(); let checksum = kernel.update_checksum(); - kernel.tear_down(); + let checksum = kernel.tear_down(); + //let t1 = now_ns(); out[i].write(KernelResult { name: kernel.name(), problem_size: kernel.default_problem_size(), reps, - total_s: (t1 - t0) as f64 * 1e-9, + launch_s: (t1 - t0) as f64 * 1e-9, + total_s: (t2 - t1) as f64 * 1e-9, checksum, }); } @@ -68,28 +91,43 @@ impl<'a> Executor<'a> { unsafe { libc::printf(c"\n".as_ptr()); libc::printf( - c"%-20s %6s %12s %16s %18s\n".as_ptr(), + c"%-20s %6s %12s %16s %16s %18s\n".as_ptr(), c"KernelBase".as_ptr(), c"Reps".as_ptr(), c"N".as_ptr(), + c"Launch time(s)".as_ptr(), + c"Sync time(s)".as_ptr(), c"Total time(s)".as_ptr(), c"Checksum".as_ptr(), ); + libc::printf( + c"%-20s %6s %12s %16s %16s %18s\n".as_ptr(), + c"----f---------------".as_ptr(), + c"------".as_ptr(), + c"------------".as_ptr(), + c"----------------".as_ptr(), + c"----------------".as_ptr(), + c"----------------".as_ptr(), + c"------------------".as_ptr(), + ); libc::printf( c"%-20s %6s %12s %16s %18s\n".as_ptr(), c"--------------------".as_ptr(), c"------".as_ptr(), c"------------".as_ptr(), c"----------------".as_ptr(), + c"----------------".as_ptr(), c"------------------".as_ptr(), ); for r in results { libc::printf( - c"%-20s %6u %12zu %16.6f %18.6f\n".as_ptr(), + c"%-20s %6u %12zu %16.6f %16.6f %18.6f\n".as_ptr(), r.name.as_bytes().as_ptr(), r.reps as libc::c_uint, r.problem_size as libc::size_t, + r.launch_s, r.total_s, + r.total_s + r.launch_s, r.checksum, ); } diff --git a/src/common/kernel_base.rs b/src/common/kernel_base.rs index fc5d037..a6a5027 100644 --- a/src/common/kernel_base.rs +++ b/src/common/kernel_base.rs @@ -6,7 +6,7 @@ pub trait KernelBase { fn setup(&mut self); fn run_kernel(&mut self); fn update_checksum(&self) -> f64; - fn tear_down(&mut self); + fn tear_down(&mut self) -> f64; } #[macro_export] diff --git a/src/common/timer.rs b/src/common/timer.rs index d02db00..1962101 100644 --- a/src/common/timer.rs +++ b/src/common/timer.rs @@ -3,7 +3,7 @@ extern crate libc; #[cfg(target_os = "linux")] #[repr(C)] -#[derive(Copy, Clone)] +#[derive(Copy, Clone, Debug)] struct Timespec { tv_sec: libc::time_t, tv_nsec: libc::c_long, diff --git a/src/common/types.rs b/src/common/types.rs index 4b619e9..3236cfe 100644 --- a/src/common/types.rs +++ b/src/common/types.rs @@ -1,7 +1,7 @@ use core::ops::{Add, AddAssign, Div, DivAssign, Mul, MulAssign, Sub, SubAssign}; #[repr(transparent)] -#[derive(Copy, Clone, Debug, Default, PartialEq, PartialOrd)] +#[derive(Copy, Debug, Clone, Default, PartialEq, PartialOrd)] pub struct A64(pub f64); impl Add for A64 { diff --git a/src/lib.rs b/src/lib.rs index 40f6055..1e8d6a4 100644 --- a/src/lib.rs +++ b/src/lib.rs @@ -1,54 +1,63 @@ #![allow(internal_features)] -#![allow(non_snake_case)] -#![allow(clippy::deref_addrof)] -#![allow(clippy::too_many_arguments)] -#![allow(clippy::missing_safety_doc)] -#![allow(unused_features)] -#![feature(abi_gpu_kernel)] -#![feature(core_float_math)] -#![feature(core_intrinsics)] -#![feature(float_algebraic)] +#![allow(linker_messages)] +#![allow(improper_ctypes)] +#![allow(improper_gpu_kernel_arg)] +#![allow(improper_ctypes_definitions)] +#![feature(stmt_expr_attributes)] +#![feature(loop_hints)] +#![feature(float_algebraic, core_float_math)] +#![cfg_attr( + target_arch = "nvptx64", + feature(stdarch_nvptx, abi_gpu_kernel, gpu_offload) +)] +#![cfg_attr( + target_arch = "amdgpu", + feature(stdarch_amdgpu, abi_gpu_kernel, gpu_offload) +)] +#![cfg_attr(any(target_arch = "nvptx64", target_arch = "amdgpu"), no_std)] +#![cfg_attr(any(target_arch = "nvptx64", target_arch = "amdgpu"), no_main)] #![feature(rustc_attrs)] -#![cfg_attr(target_arch = "amdgpu", feature(stdarch_amdgpu))] -#![cfg_attr(target_arch = "nvptx64", feature(stdarch_nvptx))] -#![no_std] +#![cfg_attr(target_os = "linux", feature(core_intrinsics, gpu_offload, offload))] +#![no_main] + +pub mod apps; +pub mod common; + +//use rust_perf; #[cfg(target_os = "linux")] extern crate libc; -#[panic_handler] -fn panic(_: &core::panic::PanicInfo) -> ! { - loop {} -} - -pub mod apps; -pub mod common; +//#[panic_handler] +//fn panic(_: &core::panic::PanicInfo) -> ! { +// loop {} +//} #[cfg(all(target_os = "linux", feature = "del_dot_vec_2d"))] -use apps::del_dot_vec_2d::DelDotVec2D; +use crate::apps::del_dot_vec_2d::DelDotVec2D; #[cfg(all(target_os = "linux", feature = "energy"))] -use apps::energy::Energy; +use crate::apps::energy::Energy; #[cfg(all(target_os = "linux", feature = "fir"))] -use apps::fir::Fir; +use crate::apps::fir::Fir; #[cfg(all(target_os = "linux", feature = "ltimes"))] -use apps::ltimes::LTimes; +use crate::apps::ltimes::LTimes; #[cfg(all(target_os = "linux", feature = "matvec_3d_stencil"))] -use apps::matvec_3d_stencil::Matvec3DStencil; +use crate::apps::matvec_3d_stencil::Matvec3DStencil; #[cfg(all(target_os = "linux", feature = "pressure"))] -use apps::pressure::Pressure; +use crate::apps::pressure::Pressure; #[cfg(all(target_os = "linux", feature = "vol3d"))] -use apps::vol3d::Vol3D; +use crate::apps::vol3d::Vol3D; +#[cfg(all(target_os = "linux", feature = "del_dot_vec_2d"))] +static mut K_DEL: DelDotVec2D = DelDotVec2D::INIT; #[cfg(all(target_os = "linux", feature = "energy"))] static mut K_ENERGY: Energy = Energy::INIT; #[cfg(all(target_os = "linux", feature = "fir"))] static mut K_FIR: Fir = Fir::INIT; -#[cfg(all(target_os = "linux", feature = "del_dot_vec_2d"))] -static mut K_DEL: DelDotVec2D = DelDotVec2D::INIT; #[cfg(all(target_os = "linux", feature = "ltimes"))] static mut K_LTIMES: LTimes = LTimes::INIT; #[cfg(all(target_os = "linux", feature = "matvec_3d_stencil"))] @@ -58,8 +67,8 @@ static mut K_PRESSURE: Pressure = Pressure::INIT; #[cfg(all(target_os = "linux", feature = "vol3d"))] static mut K_VOL3D: Vol3D = Vol3D::INIT; -#[cfg(target_os = "linux")] #[unsafe(no_mangle)] +#[cfg(target_os = "linux")] fn main() { use crate::common::executor::{Executor, KernelResult, MAX_KERNELS}; use crate::common::kernel_base::KernelBase; @@ -68,6 +77,11 @@ fn main() { let mut k_links: [Option<&mut dyn KernelBase>; MAX_KERNELS] = [const { None }; MAX_KERNELS]; let mut count = 0; + #[cfg(feature = "del_dot_vec_2d")] + { + k_links[count] = Some(unsafe { &mut *(&raw mut K_DEL) }); + count += 1; + } #[cfg(feature = "energy")] { k_links[count] = Some(unsafe { &mut *(&raw mut K_ENERGY) }); @@ -78,11 +92,6 @@ fn main() { k_links[count] = Some(unsafe { &mut *(&raw mut K_FIR) }); count += 1; } - #[cfg(feature = "del_dot_vec_2d")] - { - k_links[count] = Some(unsafe { &mut *(&raw mut K_DEL) }); - count += 1; - } #[cfg(feature = "ltimes")] { k_links[count] = Some(unsafe { &mut *(&raw mut K_LTIMES) });