diff --git a/hercules_cg/src/cpu.rs b/hercules_cg/src/cpu.rs
index 552dc3a30ea9f03813838c6c6d279f8f9e941de2..b15cf30106c76640016b1957fe9906ac01c74858 100644
--- a/hercules_cg/src/cpu.rs
+++ b/hercules_cg/src/cpu.rs
@@ -839,6 +839,8 @@ impl<'a> CPUContext<'a> {
//
// ((0 * s1 + p1) * s2 + p2) * s3 + p3 ...
let elem_size = self.codegen_type_size(elem, body)?;
+ let elem_align = get_type_alignment(&self.types, elem);
+ let aligned_elem_size = Self::round_up_to(&elem_size, elem_align, body)?;
let mut acc_offset = "0".to_string();
for (p, s) in zip(pos, dims) {
let p = self.get_value(*p, false);
@@ -848,7 +850,7 @@ impl<'a> CPUContext<'a> {
}
// Convert offset in # elements -> # bytes.
- acc_offset = Self::multiply(&acc_offset, &elem_size, body)?;
+ acc_offset = Self::multiply(&acc_offset, &aligned_elem_size, body)?;
acc_ptr = Self::gep(&acc_ptr, &acc_offset, body)?;
collect_ty = elem;
}
@@ -910,6 +912,8 @@ impl<'a> CPUContext<'a> {
// The size of an array is the size of the element multipled by
// the dynamic constant bounds.
let mut acc_size = self.codegen_type_size(elem, body)?;
+ acc_size =
+ Self::round_up_to(&acc_size, get_type_alignment(&self.types, elem), body)?;
for dc in bounds {
acc_size = Self::multiply(&acc_size, &format!("%dc{}", dc.idx()), body)?;
}
diff --git a/hercules_cg/src/gpu.rs b/hercules_cg/src/gpu.rs
index 5f2feedd28b09633e456efa4cec670d68f694fb1..a3eea2745ea6dde2929b9cd6fcc17f5c6483643f 100644
--- a/hercules_cg/src/gpu.rs
+++ b/hercules_cg/src/gpu.rs
@@ -1877,7 +1877,11 @@ namespace cg = cooperative_groups;
")".repeat(array_indices.len())
));
let element_size = self.get_size(*element_type, None);
- index_ptr.push_str(&format!(" * ({})", element_size));
+ let element_align = self.get_alignment(*element_type);
+ index_ptr.push_str(&format!(
+ " * (({} + {} - 1) / {} * {})",
+ element_size, element_align, element_align, element_align
+ ));
type_id = *element_type;
}
}
@@ -2049,7 +2053,15 @@ namespace cg = cooperative_groups;
Type::Array(element_type, extents) => {
assert!(num_fields.is_none());
let array_size = multiply_dcs(extents);
- format!("{} * {}", self.get_size(*element_type, None), array_size)
+ let elem_align = self.get_alignment(*element_type);
+ format!(
+ "(({} + {} - 1) / {} * {}) * {}",
+ self.get_size(*element_type, None),
+ elem_align,
+ elem_align,
+ elem_align,
+ array_size
+ )
}
Type::Product(fields) => {
let num_fields = num_fields.unwrap_or(fields.len());
@@ -2101,19 +2113,7 @@ namespace cg = cooperative_groups;
}
fn get_alignment(&self, type_id: TypeID) -> usize {
- match &self.types[type_id.idx()] {
- Type::Array(element_type, _) => self.get_alignment(*element_type),
- Type::Product(fields) | Type::Summation(fields) => fields
- .iter()
- .map(|field| self.get_alignment(*field))
- .max()
- .unwrap_or(0),
- Type::Boolean | Type::Integer8 | Type::UnsignedInteger8 | Type::Float8 => 1,
- Type::Integer16 | Type::UnsignedInteger16 | Type::BFloat16 => 2,
- Type::Integer32 | Type::UnsignedInteger32 | Type::Float32 => 4,
- Type::Integer64 | Type::UnsignedInteger64 | Type::Float64 => 8,
- _ => panic!("Unsupported type for alignment"),
- }
+ get_type_alignment(&self.types, type_id)
}
fn codegen_intrinsic(&self, intrinsic: &Intrinsic, ty: &Type) -> String {
diff --git a/hercules_cg/src/rt.rs b/hercules_cg/src/rt.rs
index 3db0f16f35c711372841ced574c72cc50ee007ec..8fa0c09ee512e3f2e43c5280bc3cb6947bc31dc5 100644
--- a/hercules_cg/src/rt.rs
+++ b/hercules_cg/src/rt.rs
@@ -1111,6 +1111,13 @@ impl<'a> RTContext<'a> {
//
// ((0 * s1 + p1) * s2 + p2) * s3 + p3 ...
let elem_size = self.codegen_type_size(elem);
+ let elem_align = get_type_alignment(&self.module.types, elem);
+ let aligned_elem_size = format!(
+ "(({} + {}) & !{})",
+ elem_size,
+ elem_align - 1,
+ elem_align - 1
+ );
for (p, s) in zip(pos, dims) {
let p = self.get_value(*p, bb, false);
acc_offset = format!("{} * ", acc_offset);
@@ -1119,7 +1126,7 @@ impl<'a> RTContext<'a> {
}
// Convert offset in # elements -> # bytes.
- acc_offset = format!("({} * {})", acc_offset, elem_size);
+ acc_offset = format!("({} * {})", acc_offset, aligned_elem_size);
collect_ty = elem;
}
}
@@ -1192,6 +1199,13 @@ impl<'a> RTContext<'a> {
// The size of an array is the size of the element multipled by
// the dynamic constant bounds.
let mut acc_size = self.codegen_type_size(elem);
+ let elem_align = get_type_alignment(&self.module.types, elem);
+ acc_size = format!(
+ "(({} + {}) & !{})",
+ acc_size,
+ elem_align - 1,
+ elem_align - 1
+ );
for dc in bounds {
acc_size = format!("{} * ", acc_size);
self.codegen_dynamic_constant(*dc, &mut acc_size).unwrap();
diff --git a/juno_samples/rodinia/cfd/src/setup.rs b/juno_samples/rodinia/cfd/src/setup.rs
index ad2ee9612b9959783878386d8de2bb74a270c7f6..b996f057c50c6eb404ea8a9db2d0d92dc738ac2a 100644
--- a/juno_samples/rodinia/cfd/src/setup.rs
+++ b/juno_samples/rodinia/cfd/src/setup.rs
@@ -20,11 +20,11 @@ pub const VAR_DENSITY_ENERGY: usize = VAR_MOMENTUM + NDIM;
pub const NVAR: usize = VAR_DENSITY_ENERGY + 1;
pub const deg_angle_of_attack: f32 = 0.0;
-#[repr(align(32))]
-struct Alignment([u8; 32]);
+#[repr(align(64))]
+struct Alignment([u8; 64]);
// An aligned slice is stored as a boxed slice and an offset number of elements
-// that we skip over to get the desired alignment (of 32 bytes)
+// that we skip over to get the desired alignment (of 64 bytes)
pub struct AlignedSlice<T> {
slice: Box<[T]>,
offset: usize,
@@ -37,8 +37,8 @@ where
{
pub fn of_len(len: usize) -> Self {
// The maximum number of elements that may be waisted in getting the alignment we need is
- // (32 - alignment of T) / size of T
- let extra_elements = (32 - std::mem::align_of::<T>()) / std::mem::size_of::<T>();
+ // (64 - alignment of T) / size of T
+ let extra_elements = (64 - std::mem::align_of::<T>()) / std::mem::size_of::<T>();
let slice: Box<[T]> = (0..len + extra_elements)
.map(|_| Default::default())
.collect();