GPU backend

.gitignore

@@ -4,8 +4,12 @@
 *.out
 *.ll
 *.c
+*.cu
 *.o
 *.a
 *.hrt
-.*.swp
+*.png
+*.swp
 .vscode
+*_env
+*.txt

Cargo.lock

@@ -1217,9 +1217,9 @@ checksum = "b5aba8db14291edd000dfcc4d620c7ebfb122c613afb886ca8803fa4e128a20a"
 
 [[package]]
 name = "libfuzzer-sys"
-version = "0.4.8"
+version = "0.4.9"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "9b9569d2f74e257076d8c6bfa73fb505b46b851e51ddaecc825944aa3bed17fa"
+checksum = "cf78f52d400cf2d84a3a973a78a592b4adc535739e0a5597a0da6f0c357adc75"
 dependencies = [
  "arbitrary",
  "cc",
@@ -2174,9 +2174,9 @@ dependencies = [
 
 [[package]]
 name = "toml_edit"
-version = "0.22.22"
+version = "0.22.23"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "4ae48d6208a266e853d946088ed816055e556cc6028c5e8e2b84d9fa5dd7c7f5"
+checksum = "02a8b472d1a3d7c18e2d61a489aee3453fd9031c33e4f55bd533f4a7adca1bee"
 dependencies = [
  "indexmap",
  "serde",
@@ -2444,9 +2444,9 @@ checksum = "589f6da84c646204747d1270a2a5661ea66ed1cced2631d546fdfb155959f9ec"
 
 [[package]]
 name = "winnow"
-version = "0.6.24"
+version = "0.7.0"
 source = "registry+https://github.com/rust-lang/crates.io-index"
-checksum = "c8d71a593cc5c42ad7876e2c1fda56f314f3754c084128833e64f1345ff8a03a"
+checksum = "7e49d2d35d3fad69b39b94139037ecfb4f359f08958b9c11e7315ce770462419"
 dependencies = [
  "memchr",
 ]

Cargo.toml

@@ -5,12 +5,12 @@ members = [
 	"hercules_ir",
 	"hercules_opt",
 	"hercules_rt",
-	
+
 	"juno_utils",
 	"juno_frontend",
 	"juno_scheduler",
 	"juno_build",
-	
+
 	"hercules_test/hercules_interpreter",
 	"hercules_test/hercules_tests",
 

hercules_cg/Cargo.toml

@@ -4,6 +4,9 @@ version = "0.1.0"
 authors = ["Russel Arbore <rarbore2@illinois.edu>"]
 edition = "2021"
 
+[features]
+cuda = []
+
 [dependencies]
 rand = "*"
 ordered-float = "*"

hercules_cg/src/fork_tree.rs

+use std::collections::{HashMap, HashSet};
+
+use crate::*;
+
+/*
+ * Construct a map from fork node to all control nodes (including itself) satisfying:
+ * a) domination by F
+ * b) no domination by F's join
+ * c) no domination by any other fork that's also dominated by F, where we do count self-domination
+ * Here too we include the non-fork start node, as key for all controls outside any fork.
+ */
+pub fn fork_control_map(fork_join_nesting: &HashMap<NodeID, Vec<NodeID>>) -> HashMap<NodeID, HashSet<NodeID>> {
+    let mut fork_control_map = HashMap::new();
+    for (control, forks) in fork_join_nesting {
+        let fork = forks.first().copied().unwrap_or(NodeID::new(0));
+        fork_control_map.entry(fork).or_insert_with(HashSet::new).insert(*control);
+    }
+    fork_control_map
+}
+
+/* Construct a map from each fork node F to all forks satisfying:
+ * a) domination by F
+ * b) no domination by F's join
+ * c) no domination by any other fork that's also dominated by F, where we don't count self-domination
+ * Note that the fork_tree also includes the non-fork start node, as unique root node.
+ */
+pub fn fork_tree(function: &Function, fork_join_nesting: &HashMap<NodeID, Vec<NodeID>>) -> HashMap<NodeID, HashSet<NodeID>> {
+    let mut fork_tree = HashMap::new();
+    for (control, forks) in fork_join_nesting {
+        if function.nodes[control.idx()].is_fork() {
+            fork_tree.entry(*control).or_insert_with(HashSet::new);
+            let nesting_fork = forks.get(1).copied().unwrap_or(NodeID::new(0));
+            fork_tree.entry(nesting_fork).or_insert_with(HashSet::new).insert(*control);
+        }
+    }
+    fork_tree
+}

hercules_cg/src/lib.rs

 #![feature(if_let_guard, let_chains)]
 
 pub mod cpu;
+pub mod gpu;
 pub mod rt;
 
+pub mod fork_tree;
+
 pub use crate::cpu::*;
+pub use crate::gpu::*;
 pub use crate::rt::*;
 
+pub use crate::fork_tree::*;
+
 use std::collections::BTreeMap;
 
 use hercules_ir::*;

hercules_ir/src/ir.rs

@@ -893,6 +893,14 @@ impl Type {
         }
     }
 
+    pub fn is_summation(&self) -> bool {
+        if let Type::Summation(_) = self {
+            true
+        } else {
+            false
+        }
+    }
+
     pub fn is_array(&self) -> bool {
         if let Type::Array(_, _) = self {
             true

hercules_opt/Cargo.toml

@@ -4,6 +4,9 @@ version = "0.1.0"
 authors = ["Russel Arbore <rarbore2@illinois.edu>, Aaron Councilman <aaronjc4@illinois.edu>"]
 edition = "2021"
 
+[features]
+cuda = ["hercules_cg/cuda"]
+
 [dependencies]
 ordered-float = "*"
 bitvec = "*"

hercules_rt/build.rs

@@ -11,13 +11,14 @@ fn main() {
             .status()
             .expect("PANIC: NVCC failed when building runtime. Is NVCC installed?");
         Command::new("ar")
-            .args(&["crus", "librtdefs.a", "rtdefs.o"])
             .current_dir(&Path::new(&out_dir))
+            .args(&["crus", "librtdefs.a", "rtdefs.o"])
             .status()
             .unwrap();
 
         println!("cargo::rustc-link-search=native={}", out_dir);
         println!("cargo::rustc-link-search=native=/usr/lib/x86_64-linux-gnu/");
+        println!("cargo::rustc-link-search=native=/usr/local/cuda/lib64");
         println!("cargo::rustc-link-search=native=/opt/cuda/lib/");
         println!("cargo::rustc-link-lib=static=rtdefs");
         println!("cargo::rustc-link-lib=cudart");

hercules_rt/src/lib.rs

@@ -152,6 +152,19 @@ impl<'a> HerculesCPURefMut<'a> {
 
 #[cfg(feature = "cuda")]
 impl<'a> HerculesCUDARef<'a> {
+    pub fn to_cpu_ref<'b, T>(self, dst: &'b mut [T]) -> HerculesCPURefMut<'b> {
+        unsafe {
+            let size = self.size;
+            let ptr = NonNull::new(dst.as_ptr() as *mut u8).unwrap();
+            __copy_cuda_to_cpu(ptr.as_ptr(), self.ptr.as_ptr(), size);
+            HerculesCPURefMut {
+                ptr,
+                size,
+                _phantom: PhantomData,
+            }
+        }
+    }
+
     pub unsafe fn __ptr(&self) -> *mut u8 {
         self.ptr.as_ptr()
     }
@@ -179,6 +192,19 @@ impl<'a> HerculesCUDARefMut<'a> {
         }
     }
 
+    pub fn to_cpu_ref<'b, T>(self, dst: &mut [T]) -> HerculesCPURefMut<'b> {
+        unsafe {
+            let size = self.size;
+            let ptr = NonNull::new(dst.as_ptr() as *mut u8).unwrap();
+            __copy_cuda_to_cpu(ptr.as_ptr(), self.ptr.as_ptr(), size);
+            HerculesCPURefMut {
+                ptr,
+                size,
+                _phantom: PhantomData,
+            }
+        }
+    }
+
     pub unsafe fn __ptr(&self) -> *mut u8 {
         self.ptr.as_ptr()
     }

hercules_rt/src/rtdefs.cu

@@ -7,7 +7,7 @@ extern "C" {
 		}
 		return ptr;
 	}
-	
+
 	void __cuda_dealloc(void *ptr, size_t size) {
 		(void) size;
 		cudaFree(ptr);
@@ -16,15 +16,15 @@ extern "C" {
 	void __cuda_zero_mem(void *ptr, size_t size) {
 		cudaMemset(ptr, 0, size);
 	}
-	
+
 	void __copy_cpu_to_cuda(void *dst, void *src, size_t size) {
 		cudaMemcpy(dst, src, size, cudaMemcpyHostToDevice);
 	}
-	
+
 	void __copy_cuda_to_cpu(void *dst, void *src, size_t size) {
 		cudaMemcpy(dst, src, size, cudaMemcpyDeviceToHost);
 	}
-	
+
 	void __copy_cuda_to_cuda(void *dst, void *src, size_t size) {
 		cudaMemcpy(dst, src, size, cudaMemcpyDeviceToDevice);
 	}

hercules_samples/call/Cargo.toml

@@ -4,6 +4,9 @@ version = "0.1.0"
 authors = ["Russel Arbore <rarbore2@illinois.edu>"]
 edition = "2021"
 
+[features]
+cuda = ["juno_build/cuda"]
+
 [build-dependencies]
 juno_build = { path = "../../juno_build" }
 

hercules_samples/call/build.rs

 use juno_build::JunoCompiler;
 
 fn main() {
-    JunoCompiler::new()
-        .ir_in_src("call.hir")
-        .unwrap()
-        .build()
-        .unwrap();
+    #[cfg(not(feature = "cuda"))]
+    {
+        JunoCompiler::new()
+            .ir_in_src("call.hir")
+            .unwrap()
+            .build()
+            .unwrap();
+    }
+    #[cfg(feature = "cuda")]
+    {
+        JunoCompiler::new()
+            .ir_in_src("call.hir")
+            .unwrap()
+            .schedule_in_src("gpu.sch")
+            .unwrap()
+            .build()
+            .unwrap();
+    }
 }

hercules_samples/call/src/gpu.sch

+gvn(*);
+phi-elim(*);
+dce(*);
+
+let out = auto-outline(*);
+gpu(out.add);
+
+ip-sroa(*);
+sroa(*);
+dce(*);
+gvn(*);
+phi-elim(*);
+dce(*);
+
+infer-schedules(*);
+
+gcm(*);

hercules_samples/ccp/Cargo.toml

@@ -4,8 +4,8 @@ version = "0.1.0"
 authors = ["Russel Arbore <rarbore2@illinois.edu>"]
 edition = "2021"
 
-[build-dependencies]
-juno_build = { path = "../../juno_build" }
+[features]
+cuda = ["juno_build/cuda"]
 
 [dependencies]
 juno_build = { path = "../../juno_build" }
@@ -13,3 +13,6 @@ hercules_rt = { path = "../../hercules_rt" }
 rand = "*"
 async-std = "*"
 with_builtin_macros = "0.1.0"
+
+[build-dependencies]
+juno_build = { path = "../../juno_build" }

hercules_samples/ccp/build.rs

 use juno_build::JunoCompiler;
 
 fn main() {
-    JunoCompiler::new()
-        .ir_in_src("ccp.hir")
-        .unwrap()
-        .build()
-        .unwrap();
+    #[cfg(not(feature = "cuda"))]
+    {
+        JunoCompiler::new()
+            .ir_in_src("ccp.hir")
+            .unwrap()
+            .build()
+            .unwrap();
+    }
+    #[cfg(feature = "cuda")]
+    {
+        JunoCompiler::new()
+            .ir_in_src("ccp.hir")
+            .unwrap()
+            .schedule_in_src("gpu.sch")
+            .unwrap()
+            .build()
+            .unwrap();
+    }
 }

hercules_samples/ccp/src/gpu.sch

+gvn(*);
+phi-elim(*);
+dce(*);
+
+let out = auto-outline(*);
+gpu(out.tricky);
+
+ip-sroa(*);
+sroa(*);
+dce(*);
+gvn(*);
+phi-elim(*);
+dce(*);
+
+infer-schedules(*);
+
+gcm(*);

hercules_samples/dot/Cargo.toml

@@ -5,7 +5,7 @@ authors = ["Russel Arbore <rarbore2@illinois.edu>"]
 edition = "2021"
 
 [features]
-cuda = ["hercules_rt/cuda"]
+cuda = ["juno_build/cuda", "hercules_rt/cuda"]
 
 [build-dependencies]
 juno_build = { path = "../../juno_build" }

hercules_samples/dot/build.rs

@@ -4,8 +4,7 @@ fn main() {
     JunoCompiler::new()
         .ir_in_src("dot.hir")
         .unwrap()
-        //.schedule_in_src(if cfg!(feature = "cuda") { "gpu.sch" } else { "cpu.sch" })
-        .schedule_in_src("cpu.sch")
+        .schedule_in_src(if cfg!(feature = "cuda") { "gpu.sch" } else { "cpu.sch" })
         .unwrap()
         .build()
         .unwrap();