From 8070fc0baa2291e1635bb347fdd032dffa129dd5 Mon Sep 17 00:00:00 2001
From: Russel Arbore <rarbore2@illinois.edu>
Date: Fri, 7 Mar 2025 18:27:23 -0600
Subject: [PATCH] don't emit sync shared mem if no syncs
---
hercules_cg/src/gpu.rs | 38 +++++++++++++++++++++++++-------------
1 file changed, 25 insertions(+), 13 deletions(-)
diff --git a/hercules_cg/src/gpu.rs b/hercules_cg/src/gpu.rs
index 4069cb02..c07ed4e5 100644
--- a/hercules_cg/src/gpu.rs
+++ b/hercules_cg/src/gpu.rs
@@ -1,6 +1,7 @@
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
--
GitLab