search

bheisler / RustaCUDA

Rusty wrapper for the CUDA Driver API
Apache License 2.0
765 stars 58 forks source link

Add cooperative kernel launch API #62

Open LutzCle opened 2 years ago

LutzCle commented 2 years ago

Pascal and newer devices support cooperative groups. These groups enable kernels to, e.g., globally synchronize the grid without terminating the kernel.

A special function cuLaunchCooperativeKernel was added to CUDA 9.0 to launch a cooperative group. This PR adds this launch API to RustaCUDA.

Example:

unsafe { launch_cooperative!( kernel<<<grid, block>>>() ).unwrap() };

See the CUDA documentation and this Nvidia Developer blog post for more information.