Add bridging PTX code for cooperative_groups API

This appears to be a working solution. The PTX should be fairly portable
as well. Overall, this system will work as follows:
- Wrap a `cc::Builder` to compile the bridging code
  on demand as part of a `sys` crate. This will use nvcc under the hood
  and we will pass along any other needed flags.
- Folks that need the bridging code will then use cust::link::Linker to link
  the bridging PTX code with their PTX.

All of the above is tested and working, however it currently deadlocks
invocations because we have not exposed the cooperative launch
interface. This should be quite simple though, given that the generated
cuda bindgen code already has this in place.

LOTS TO DO STILL!
- Update the new cuda_std_cg crate to use `cc` to compile the C++
  bridging code on demand, and produce the PTX path as output.
- Remove the Justfile. I was only using it for POC testing.
- MAJOR: update cust to expose the `cuLaunchCooperativeKernel` in a nice
  interface.

Author: thedodd

Cargo.toml

@@ -14,7 +14,8 @@ members = [
 ]
 
 exclude = [
-    "crates/optix/examples/common"
+    "crates/optix/examples/common",
+    "crates/cuda_std_cg",
 ]
 
 [profile.dev.package.rustc_codegen_nvvm]

Justfile

@@ -0,0 +1,7 @@
+build_cuda_std_cg:
+    #!/usr/bin/env bash
+    set -euxo pipefail
+    nvcc --ptx -arch=sm_75 \
+        -I crates/cuda_std_cg/src -I${CUDA_ROOT}/include \
+        --device-c crates/cuda_std_cg/src/cg_bridge.cu \
+        -o crates/cuda_std_cg/cg_bridge.ptx

crates/cuda_std/src/cg.rs

@@ -0,0 +1,72 @@
+use crate::gpu_only;
+
+mod ffi {
+    use core::ffi::c_void;
+
+    pub type GridGroup = *mut c_void;
+    extern "C" {
+        pub(super) fn this_grid() -> GridGroup;
+        pub(super) fn GridGroup_destroy(gg: GridGroup);
+        pub(super) fn GridGroup_is_valid(gg: GridGroup) -> bool;
+        pub(super) fn GridGroup_sync(gg: GridGroup);
+        pub(super) fn GridGroup_size(gg: GridGroup) -> u64;
+        pub(super) fn GridGroup_thread_rank(gg: GridGroup) -> u64;
+        pub(super) fn GridGroup_num_threads(gg: GridGroup) -> u64;
+        pub(super) fn GridGroup_num_blocks(gg: GridGroup) -> u64;
+        pub(super) fn GridGroup_block_rank(gg: GridGroup) -> u64;
+        // dim3 GridGroup_group_dim(); // TODO: impl these.
+        // dim3 GridGroup_dim_blocks(); // TODO: impl these.
+        // dim3 GridGroup_block_index(); // TODO: impl these.
+    }
+}
+
+pub struct GridGroup(ffi::GridGroup);
+
+impl Drop for GridGroup {
+    fn drop(&mut self) {
+        unsafe { ffi::GridGroup_destroy(self.0) }
+    }
+}
+
+impl GridGroup {
+    #[gpu_only]
+    pub fn this_grid() -> Self {
+        let ptr = unsafe { ffi::this_grid() };
+        GridGroup(ptr)
+    }
+
+    #[gpu_only]
+    pub fn is_valid(&mut self) -> bool {
+        unsafe { ffi::GridGroup_is_valid(self.0) }
+    }
+
+    #[gpu_only]
+    pub fn sync(&mut self) {
+        unsafe { ffi::GridGroup_sync(self.0) }
+    }
+
+    #[gpu_only]
+    pub fn size(&mut self) -> u64 {
+        unsafe { ffi::GridGroup_size(self.0) }
+    }
+
+    #[gpu_only]
+    pub fn thread_rank(&mut self) -> u64 {
+        unsafe { ffi::GridGroup_thread_rank(self.0) }
+    }
+
+    #[gpu_only]
+    pub fn num_threads(&mut self) -> u64 {
+        unsafe { ffi::GridGroup_num_threads(self.0) }
+    }
+
+    #[gpu_only]
+    pub fn num_blocks(&mut self) -> u64 {
+        unsafe { ffi::GridGroup_num_blocks(self.0) }
+    }
+
+    #[gpu_only]
+    pub fn block_rank(&mut self) -> u64 {
+        unsafe { ffi::GridGroup_block_rank(self.0) }
+    }
+}

crates/cuda_std/src/lib.rs

@@ -46,6 +46,7 @@ pub mod misc;
 // pub mod rt;
 pub mod atomic;
 pub mod cfg;
+pub mod cg;
 pub mod ptr;
 pub mod shared;
 pub mod thread;