Sync optimizations

• Optimize emitting of synchronization in GPU backend (in particular, avoid it sometimes).

hercules_cg/src/gpu.rs

 extern crate bitvec;
 extern crate hercules_ir;
 
+use std::cell::RefCell;
 use std::collections::{BTreeMap, HashMap, HashSet};
 use std::fmt::{Error, Write};
 
@@ -191,6 +192,7 @@ pub fn gpu_codegen<W: Write>(
         control_data_phi_map,
         return_parameters,
         kernel_params,
+        generated_sync: RefCell::new(false),
     };
     ctx.codegen_function(w)
 }
@@ -221,6 +223,7 @@ struct GPUContext<'a> {
     control_data_phi_map: HashMap<NodeID, Vec<(NodeID, NodeID)>>,
     return_parameters: Vec<Option<usize>>,
     kernel_params: &'a GPUKernelParams,
+    generated_sync: RefCell<bool>,
 }
 
 /*
@@ -277,8 +280,6 @@ impl GPUContext<'_> {
         let mut dynamic_shared_offset = "0".to_string();
         self.codegen_dynamic_constants(&mut top)?;
         self.codegen_declare_data(&mut top)?;
-        self.codegen_helpers(&mut top)?;
-        write!(w, "{}", top)?;
 
         // Setup for CUDA's "goto" for control flow between basic blocks.
         let mut gotos: BTreeMap<_, _> = (0..self.function.nodes.len())
@@ -327,6 +328,8 @@ impl GPUContext<'_> {
         };
 
         // Emit all GPU kernel code from previous steps
+        self.codegen_helpers(&mut top, *self.generated_sync.borrow())?;
+        write!(w, "{}", top)?;
         self.codegen_goto_start(&mut thread_block_tiles)?;
         write!(w, "{}", thread_block_tiles)?;
         let mut kernel_body = String::new();
@@ -575,16 +578,25 @@ namespace cg = cooperative_groups;
      * are from CUDA's cooperative groups API and are used specifically for reads
      * and writes.
      */
-    fn codegen_helpers(&self, w: &mut String) -> Result<(), Error> {
-        write!(
-            w,
-            "\t__shared__ cge::block_tile_memory<1024> block_sync_shared;\n"
-        )?;
+    fn codegen_helpers(&self, w: &mut String, need_sync_shared: bool) -> Result<(), Error> {
+        if need_sync_shared {
+            write!(
+                w,
+                "\t__shared__ cge::block_tile_memory<1024> block_sync_shared;\n"
+            )?;
+        }
         write!(w, "\tcg::grid_group grid = cg::this_grid();\n")?;
-        write!(
-            w,
-            "\tcg::thread_block block = cge::this_thread_block(block_sync_shared);\n"
-        )?;
+        if need_sync_shared {
+            write!(
+                w,
+                "\tcg::thread_block block = cge::this_thread_block(block_sync_shared);\n"
+            )?;
+        } else {
+            write!(
+                w,
+                "\tcg::thread_block block = cg::this_thread_block();\n"
+            )?;
+        }
         Ok(())
     }
 
@@ -1344,7 +1356,7 @@ namespace cg = cooperative_groups;
                     write!(w, "{}\t*({} + i) = 0;\n", tabs, define_variable)?;
                     write!(w, "{}}}\n", tabs)?;
                     write!(w, "{}{}.sync();\n", tabs, cg_tile)?;
-                    //write!(w, "__syncthreads\n")?;
+                    *self.generated_sync.borrow_mut() = true;
                 }
             }
             // Dynamic constants emitted at top
@@ -1806,7 +1818,7 @@ namespace cg = cooperative_groups;
                     let fork = self.join_fork_map.get(&id).unwrap();
                     let cg_tile_available = self.get_cg_tile(*fork, CGType::Available);
                     write!(w_term, "\t{}.sync();\n", cg_tile_available)?;
-                    //write!(w_term, "\t__syncthreads;\n")?;
+                    *self.generated_sync.borrow_mut() = true;
                 }
                 // If the Fork was parallelized, each thread or UsedPerId tile of
                 // threads only runs one ThreadID, so we can jump straight to the