Skip to content

Enable code for dynamic parallelism #96

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
wants to merge 3 commits into
base: main
Choose a base branch
from
Open

Enable code for dynamic parallelism #96

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
3fd5949
Enable code for dynamic parallelism
thedodd Nov 15, 2022
7d622ca
WIP: working solution, but a lot of cruft. Clean up and unwind the
thedodd Nov 20, 2022
95f1066
WIP: pop this and continue experimentation
thedodd Mar 30, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
Prev Previous commit
WIP: pop this and continue experimentation
  • Loading branch information
@thedodd
thedodd committed Mar 30, 2023
commit 95f1066d5c07e1a99e8f71db32463a005ca98f3a
78 changes: 7 additions & 71 deletions crates/cuda_std/src/rt/mod.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -35,11 +35,6 @@ pub struct Stream {
}

impl Stream {
// /// Creates a new stream with flags.
// pub fn new(flags: StreamFlags) -> Self {
// Self {}
// }

/// Creates a new stream with flags.
pub fn new(flags: StreamFlags) -> CudaResult<Self> {
let mut stream = MaybeUninit::uninit();
Expand All @@ -52,11 +47,10 @@ impl Stream {
}
}

// #[doc(hidden)]
// pub fn launch(&self, param_buf: *mut c_void) -> CudaResult<()> {
// unsafe { cuda::cudaLaunchDeviceV2(param_buf, core::ptr::null_mut()).to_result() }
// // unsafe { cuda::cudaLaunchDeviceV2(param_buf, self.raw).to_result() }
// }
#[doc(hidden)]
pub unsafe fn launch(&self, param_buf: *mut c_void) -> CudaResult<()> {
cuda::cudaLaunchDeviceV2(param_buf, self.raw).to_result()
}
}

impl Drop for Stream {
Expand All @@ -69,8 +63,7 @@ impl Drop for Stream {

#[macro_export]
macro_rules! launch {
// ($func:ident<<<$grid_dim:expr, $block_dim:expr, $smem_size:expr, $stream:ident>>>($($param:expr),* $(,)?)) => {{
($func:ident<<<$grid_dim:expr, $block_dim:expr, ($smem_size:expr)>>>($($param:expr),* $(,)?)) => {{
($func:ident<<<$grid_dim:expr, $block_dim:expr, $smem_size:expr, $stream:ident>>>($($param:expr),* $(,)?)) => {{
use $crate::rt::ToResult;
use $crate::float::GpuFloat;
let grid_dim = $crate::rt::GridSize::from($grid_dim);
Expand Down Expand Up @@ -108,69 +101,12 @@ macro_rules! launch {
offset += size;
)*
if false {
// Ensure function call compatibility at compile time.
$func($($param),*);
}
// unsafe {
// let mut offset = 0;
// $(
// let param = $param;
// let size = ::core::mem::size_of_val(&param);
// let mut buf_idx = (offset as f32 / size as f32).ceil() as usize + 1;
// offset = buf_idx * size;
// let ptr = &param as *const _ as *const u8;
// let dst = buf.add(offset);
// ::core::ptr::copy_nonoverlapping(&param as *const _ as *const u8, dst, size);
// )*
// }
// if false {
// $func($($param),*);
// }

// Launch the kernel.
$crate::rt::sys::cudaLaunchDeviceV2(buf as *mut ::core::ffi::c_void, ::core::ptr::null_mut() as *mut _)

// let mut buf = $crate::rt::sys::cudaGetParameterBuffer(alignment, size) as *mut u8;

// // Populate the buffer with given arguments.
// let mut offset = 0;
// $(
// let param = $param;
// let size = ::core::mem::size_of_val(&param);
// let buf_bytes_ptr = (buf as *mut u8).add(offset);
// ::core::ptr::copy_nonoverlapping($param as *const _, buf_bytes_ptr.into(), size);
// offset += size;
// )*

// let mut offset = 0;
// $(
// let param = $param;
// let size = ::core::mem::size_of_val(&param);
// let mut buf_idx = (offset as f32 / size as f32).ceil() as usize + 1;
// offset = buf_idx * size;
// let ptr = &param as *const _ as *const u8;
// let dst = buf.add(offset);
// ::core::ptr::copy_nonoverlapping(&param as *const _ as *const u8, dst, size);
// )*

// // Launch the kernel.
// let fptr = $func as *const ();
// $crate::rt::sys::cudaLaunchDevice(
// fptr as *const ::core::ffi::c_void,
// buf as *mut ::core::ffi::c_void,
// $crate::rt::sys::dim3 {
// x: grid_dim.x,
// y: grid_dim.y,
// z: grid_dim.z
// },
// $crate::rt::sys::dim3 {
// x: block_dim.x,
// y: block_dim.y,
// z: block_dim.z
// },
// $smem_size,
// ::core::ptr::null_mut() as *mut _,
// // $stream.raw,
// )
$stream.launch(buf)
}};
}

Expand Down
1 change: 1 addition & 0 deletions crates/cust/Cargo.toml
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -22,6 +22,7 @@ mint = { version = "^0.5", optional = true }
num-complex = { version = "0.4", optional = true }
vek = { version = "0.15.1", optional = true, default-features = false }
bytemuck = { version = "1.7.3", optional = true }
find_cuda_helper = { path = "../find_cuda_helper", version = "0.2" }

[features]
default= ["bytemuck"]
Expand Down
18 changes: 12 additions & 6 deletions crates/cust/src/link.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3,9 +3,9 @@
use std::mem::MaybeUninit;
use std::ptr::null_mut;

use crate::error::{CudaError, CudaResult, ToResult};
use crate::sys as cuda;

use crate::error::{CudaResult, ToResult};
use find_cuda_helper::find_lib_cudadevrt;

static UNNAMED: &str = "\0";

Expand All @@ -25,6 +25,12 @@ impl Linker {
// Therefore we use box to alloc the memory for us, then into_raw it so we now have ownership
// of the memory (and dont have any aliasing requirements attached either).

// // Just take advantage of C memory model and just pass individual elements, as there is only 1.
// let num_options: u32 = 1;
// let opt = &mut cuda::CUjit_option::CU_JIT_TARGET as *mut _;
// let mut opt_val =
// &mut cuda::CUjit_target::CU_TARGET_COMPUTE_75 as *mut _ as *mut ::std::os::raw::c_void;

unsafe {
let mut raw = MaybeUninit::uninit();
cuda::cuLinkCreate_v2(0, null_mut(), null_mut(), raw.as_mut_ptr()).to_result()?;
Expand Down Expand Up @@ -116,15 +122,15 @@ impl Linker {

/// Link device runtime lib.
pub fn add_libcudadevrt(&mut self) -> CudaResult<()> {
let mut bytes = std::fs::read("/usr/local/cuda-11/lib64/libcudadevrt.a")
.expect("could not read libcudadevrt.a");
let path = find_lib_cudadevrt().ok_or_else(|| CudaError::FileNotFound)?;
let mut bytes = std::fs::read(path)
// TODO: don't panic, update the result type instead.
.expect("error linking libcudadevrt.a");

unsafe {
cuda::cuLinkAddData_v2(
self.raw,
cuda::CUjitInputType::CU_JIT_INPUT_LIBRARY,
// cuda_sys wants *mut but from the API docs we know we retain ownership so
// this cast is sound.
bytes.as_mut_ptr() as *mut _,
bytes.len(),
UNNAMED.as_ptr().cast(),
Expand Down
10 changes: 10 additions & 0 deletions crates/find_cuda_helper/src/lib.rs
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -150,6 +150,16 @@ pub fn find_cuda_lib_dirs() -> Vec<PathBuf> {
valid_paths
}

/// Find the location of `libcudadevrt.a`.
pub fn find_lib_cudadevrt() -> Option<PathBuf> {
let root = find_cuda_root()?;
let lib = root.join("lib64").join("libcudadevrt.a");
if lib.is_file() {
return Some(lib);
}
None
}

#[cfg(target_os = "windows")]
pub fn find_optix_root() -> Option<PathBuf> {
// the optix SDK installer sets OPTIX_ROOT_DIR whenever it installs.
Expand Down