(ml5717) Implemented CUDA specific Event Buffer

Author: juntyr

Cargo.toml

@@ -11,6 +11,7 @@ members = [
     "necsim-core",
     "necsim-impls-std",
     "necsim-impls-no-std",
+    "necsim-impls-cuda",
 
     "necsim-classical",
     "necsim-gillespie",

TODO

@@ -1,4 +1,5 @@
-- implement event transform from GPU to CPU
+- implement event transfer from GPU to CPU
+- implement simulation domain slicing so that event buffer fits into GPU RAM
 - implement event deduplication (global vs local? - needs to be global for speciation)
 - implement simple stop condition for time sliced kernel launching
 - implement declarative event filtering which can be put onto the GPU as well

necsim-core/src/event.rs

@@ -3,11 +3,17 @@ use core::marker::PhantomData;
 use crate::cogs::{Habitat, LineageReference};
 use crate::landscape::IndexedLocation;
 
+#[cfg(feature = "cuda")]
+use rustacuda_core::DeviceCopy;
+
+#[cfg(feature = "cuda")]
+use rust_cuda::common::RustToCuda;
+
 pub struct Event<H: Habitat, R: LineageReference<H>> {
     time: f64,
     lineage_reference: R,
     r#type: EventType<H, R>,
-    _marker: PhantomData<H>,
+    marker: PhantomData<H>,
 }
 
 impl<H: Habitat, R: LineageReference<H>> Event<H, R> {
@@ -21,7 +27,7 @@ impl<H: Habitat, R: LineageReference<H>> Event<H, R> {
             time,
             lineage_reference,
             r#type,
-            _marker: PhantomData::<H>,
+            marker: PhantomData::<H>,
         }
     }
 
@@ -49,10 +55,40 @@ pub enum EventType<H: Habitat, R: LineageReference<H>> {
         origin: IndexedLocation,
         target: IndexedLocation,
         coalescence: Option<R>,
-        _marker: PhantomData<H>,
+        marker: PhantomData<H>,
     },
 }
 
+impl<H: Habitat, R: LineageReference<H>> Clone for Event<H, R> {
+    fn clone(&self) -> Self {
+        Self {
+            time: self.time,
+            lineage_reference: self.lineage_reference.clone(),
+            r#type: self.r#type.clone(),
+            marker: self.marker.clone(),
+        }
+    }
+}
+
+impl<H: Habitat, R: LineageReference<H>> Clone for EventType<H, R> {
+    fn clone(&self) -> Self {
+        match self {
+            EventType::Speciation => EventType::Speciation,
+            EventType::Dispersal {
+                origin,
+                target,
+                coalescence,
+                marker,
+            } => EventType::Dispersal {
+                origin: origin.clone(),
+                target: target.clone(),
+                coalescence: coalescence.clone(),
+                marker: marker.clone(),
+            },
+        }
+    }
+}
+
 impl<H: Habitat, R: LineageReference<H>> core::fmt::Debug for Event<H, R> {
     fn fmt(&self, f: &mut core::fmt::Formatter<'_>) -> core::fmt::Result {
         f.debug_struct("Event")
@@ -62,7 +98,19 @@ impl<H: Habitat, R: LineageReference<H>> core::fmt::Debug for Event<H, R> {
             )
             .field("lineage_reference", &self.lineage_reference)
             .field("type", &self.r#type)
-            .field("_marker", &format_args!("PhantomData"))
+            .field("marker", &format_args!("PhantomData"))
             .finish()
     }
 }
+
+#[cfg(feature = "cuda")]
+unsafe impl<H: Habitat + RustToCuda, R: LineageReference<H> + DeviceCopy> DeviceCopy
+    for Event<H, R>
+{
+}
+
+#[cfg(feature = "cuda")]
+unsafe impl<H: Habitat + RustToCuda, R: LineageReference<H> + DeviceCopy> DeviceCopy
+    for EventType<H, R>
+{
+}

necsim-cuda/Cargo.toml

@@ -10,6 +10,7 @@ edition = "2018"
 necsim-core = { path = "../necsim-core", features = ["cuda"] }
 necsim-impls-no-std = { path = "../necsim-impls-no-std", features = ["cuda"] }
 necsim-impls-std = { path = "../necsim-impls-std" }
+necsim-impls-cuda = { path = "../necsim-impls-cuda" }
 array2d = { path = "../array2d-no-std" }
 anyhow = "1.0"
 contracts = "0.6.0"

necsim-cuda/kernel/Cargo.toml

@@ -11,4 +11,5 @@ ryu = "1.0.5"
 necsim-core = { path = "../../necsim-core", features = ["cuda"] }
 rust-cuda = { path = "../../rust-cuda", features = [] }
 necsim-impls-no-std = { path = "../../necsim-impls-no-std", features = ["cuda"] }
+necsim-impls-cuda = { path = "../../necsim-impls-cuda" }
 rustacuda_core = "0.1.2"

necsim-cuda/kernel/src/lib.rs

@@ -43,19 +43,21 @@ use necsim_core::cogs::{
     ActiveLineageSampler, CoalescenceSampler, DispersalSampler, EventSampler,
     HabitatToU64Injection, IncoherentLineageStore, LineageReference, PrimeableRng,
 };
-use necsim_core::reporter::NullReporter;
 use necsim_core::simulation::Simulation;
 use rust_cuda::common::RustToCuda;
 use rust_cuda::device::BorrowFromRust;
 use rustacuda_core::DeviceCopy;
 
-use necsim_impls_no_std::cogs::rng::cuda::CudaRng;
+use necsim_impls_cuda::cogs::rng::CudaRng;
+use necsim_impls_cuda::event_buffer::common::EventBufferCudaRepresentation;
+use necsim_impls_cuda::event_buffer::device::EventBufferDevice;
 
 #[no_mangle]
 /// # Safety
 /// This CUDA kernel is unsafe as it is called with raw pointers
 pub unsafe extern "ptx-kernel" fn simulate(
     simulation_c_ptr: *mut core::ffi::c_void,
+    event_buffer_c_ptr: *mut core::ffi::c_void,
     max_steps: usize,
 ) {
     use necsim_impls_no_std::cogs::active_lineage_sampler::independent::IndependentActiveLineageSampler as ActiveLineageSampler;
@@ -79,6 +81,7 @@ pub unsafe extern "ptx-kernel" fn simulate(
                 EventSampler<_, _, _, _, _>,
                 ActiveLineageSampler<_, _, _, _, _>,
             > as RustToCuda>::CudaRepresentation,
+        event_buffer_c_ptr as *mut EventBufferCudaRepresentation<Habitat, LineageReference>,
         max_steps,
     )
 }
@@ -94,20 +97,21 @@ unsafe fn simulate_generic<
     A: ActiveLineageSampler<H, G, D, R, S, C, E> + RustToCuda,
 >(
     simulation_ptr: *mut <Simulation<H, G, D, R, S, C, E, A> as RustToCuda>::CudaRepresentation,
+    event_buffer_ptr: *mut EventBufferCudaRepresentation<H, R>,
     max_steps: usize,
 ) {
     Simulation::with_borrow_from_rust_mut(simulation_ptr, |simulation| {
-        let mut reporter = NullReporter;
+        EventBufferDevice::with_borrow_from_rust_mut(event_buffer_ptr, |event_buffer_reporter| {
+            let (time, steps) = simulation.simulate_incremental(max_steps, event_buffer_reporter);
 
-        let (time, steps) = simulation.simulate_incremental(max_steps, &mut reporter);
-
-        if utils::thread_idx().as_id(&utils::block_dim()) == 0 {
-            println!(
-                "index = {}, time = {:?}, steps = {}",
-                utils::index(),
-                F64(time),
-                steps
-            );
-        }
+            if utils::thread_idx().as_id(&utils::block_dim()) == 0 {
+                println!(
+                    "index = {}, time = {:?}, steps = {}",
+                    utils::index(),
+                    F64(time),
+                    steps
+                );
+            }
+        })
     })
 }

necsim-cuda/src/lib.rs

@@ -21,14 +21,15 @@ use necsim_core::cogs::{LineageStore, PrimeableRng};
 use necsim_core::reporter::Reporter;
 use necsim_core::simulation::Simulation;
 
+use necsim_impls_cuda::cogs::rng::CudaRng;
+use necsim_impls_cuda::event_buffer::host::EventBufferHost;
 use necsim_impls_no_std::cogs::active_lineage_sampler::independent::IndependentActiveLineageSampler;
 use necsim_impls_no_std::cogs::coalescence_sampler::independent::IndependentCoalescenceSampler;
 use necsim_impls_no_std::cogs::dispersal_sampler::in_memory::packed_alias::InMemoryPackedAliasDispersalSampler;
 use necsim_impls_no_std::cogs::event_sampler::independent::IndependentEventSampler;
 use necsim_impls_no_std::cogs::habitat::in_memory::InMemoryHabitat;
 use necsim_impls_no_std::cogs::lineage_reference::in_memory::InMemoryLineageReference;
 use necsim_impls_no_std::cogs::lineage_store::incoherent::in_memory::IncoherentInMemoryLineageStore;
-use necsim_impls_no_std::cogs::rng::cuda::CudaRng;
 use necsim_impls_std::cogs::dispersal_sampler::in_memory::InMemoryDispersalSampler;
 
 macro_rules! with_cuda {
@@ -155,8 +156,10 @@ impl CudaSimulation {
         speciation_probability_per_generation: f64,
         sample_percentage: f64,
         rng: G,
-        _reporter: &mut impl Reporter<InMemoryHabitat, InMemoryLineageReference>,
+        reporter: &mut impl Reporter<InMemoryHabitat, InMemoryLineageReference>,
     ) -> Result<(f64, usize)> {
+        const SIMULATION_STEP_SLICE: usize = 1_000_usize;
+
         let habitat = InMemoryHabitat::new(habitat.clone());
         let dispersal_sampler = InMemoryPackedAliasDispersalSampler::new(dispersal, &habitat)?;
         let lineage_store = IncoherentInMemoryLineageStore::new(sample_percentage, &habitat);
@@ -187,6 +190,9 @@ impl CudaSimulation {
                 + (total_individuals % cuda_block_length > 0) as u32
         });
 
+        let mut event_buffer: EventBufferHost<InMemoryHabitat, InMemoryLineageReference> =
+            EventBufferHost::new(&cuda_block_size, &cuda_grid_size, SIMULATION_STEP_SLICE)?;
+
         //let (time, steps) = simulation.simulate(rng, reporter);
 
         let module_data = CString::new(include_str!(env!("KERNEL_PTX_PATH"))).unwrap();
@@ -216,16 +222,27 @@ impl CudaSimulation {
             print_kernel_function_attributes(&simulate_kernel);
 
             if let Err(err) = simulation.lend_to_cuda_mut(|simulation_mut_ptr| {
+                let block_index_range = 0..(cuda_grid_size.x * cuda_grid_size.y * cuda_grid_size.z);
+
                 // Launching kernels is unsafe since Rust can't enforce safety - think of kernel launches
                 // as a foreign-function call. In this case, it is - this kernel is written in CUDA C.
                 unsafe {
                     launch!(simulate_kernel<<<cuda_grid_size, cuda_block_size, 0, stream>>>(
                         simulation_mut_ptr,
-                        1_000_usize // max steps on GPU
+                        event_buffer.get_mut_cuda_ptr(),
+                        SIMULATION_STEP_SLICE
                     ))?;
                 }
 
-                stream.synchronize()
+                stream.synchronize()?;
+
+                for block_index in block_index_range {
+                    event_buffer.with_fetched_events_for_block(block_index as usize, |events| {
+                        events.iter().for_each(|event| reporter.report_event(event))
+                    })?
+                }
+
+                Ok(())
             }) {
                 eprintln!("Running kernel failed with {:#?}!", err);
             }

necsim-impls-cuda/Cargo.toml

@@ -0,0 +1,21 @@
+[package]
+name = "necsim-impls-cuda"
+version = "0.1.0"
+authors = ["Moritz Langenstein <ml5717@ic.ac.uk>"]
+edition = "2018"
+
+# See more keys and their definitions at https://doc.rust-lang.org/cargo/reference/manifest.html
+
+[dependencies]
+necsim-core = { path = "../necsim-core", features = ["cuda"] }
+contracts = "0.6.0"
+
+rust-cuda-derive = { path = "../rust-cuda/rust-cuda-derive" }
+rustacuda_core = { version = "0.1.2" }
+
+[target.'cfg(target_os = "cuda")'.dependencies]
+rust-cuda = { path = "../rust-cuda", features = [] }
+
+[target.'cfg(not(target_os = "cuda"))'.dependencies]
+rust-cuda = { path = "../rust-cuda", features = ["host"] }
+rustacuda = { version = "0.1.2" }

necsim-impls-cuda/src/cogs/mod.rs

@@ -0,0 +1 @@
+pub mod rng;