Fix type layout

• Make sure all elements in nested array are aligned.
• Update AlignedSlice in CFD.

hercules_cg/src/gpu.rs

@@ -1879,10 +1879,8 @@ namespace cg = cooperative_groups;
                     let element_size = self.get_size(*element_type, None);
                     let element_align = self.get_alignment(*element_type);
                     index_ptr.push_str(&format!(
-                        " * (({} + {}) & ~{})",
-                        element_size,
-                        element_align - 1,
-                        element_align - 1
+                        " * (({} + {} - 1) / {} * {})",
+                        element_size, element_align, element_align, element_align
                     ));
                     type_id = *element_type;
                 }
@@ -1934,10 +1932,8 @@ namespace cg = cooperative_groups;
                     let alignment = self.get_alignment(*type_id);
                     let size = self.get_size(*type_id, None);
                     *dynamic_shared_offset = format!(
-                        "(({} + {}) & ~{})",
-                        dynamic_shared_offset,
-                        alignment - 1,
-                        alignment - 1,
+                        "(({} + {} - 1) / {}) * {}",
+                        dynamic_shared_offset, alignment, alignment, alignment
                     );
                     write!(
                         w,
@@ -1985,10 +1981,8 @@ namespace cg = cooperative_groups;
                     let alignment = self.get_alignment(*type_id);
                     let size = self.get_size(*type_id, None);
                     *dynamic_shared_offset = format!(
-                        "(({} + {}) & ~{})",
-                        dynamic_shared_offset,
-                        alignment - 1,
-                        alignment - 1
+                        "(({} + {} - 1) / {}) * {}",
+                        dynamic_shared_offset, alignment, alignment, alignment
                     );
                     write!(
                         w,
@@ -2029,10 +2023,8 @@ namespace cg = cooperative_groups;
                 let alignment = self.get_alignment(*type_id);
                 let size = self.get_size(*type_id, None);
                 *dynamic_shared_offset = format!(
-                    "(({} + {}) & ~{})",
-                    dynamic_shared_offset,
-                    alignment - 1,
-                    alignment - 1,
+                    "(({} + {} - 1) / {}) * {}",
+                    dynamic_shared_offset, alignment, alignment, alignment
                 );
                 write!(
                     w,
@@ -2063,10 +2055,11 @@ namespace cg = cooperative_groups;
                 let array_size = multiply_dcs(extents);
                 let elem_align = self.get_alignment(type_id);
                 format!(
-                    "(({} + {}) & ~{}) * {}",
+                    "(({} + {} - 1) / {} * {}) * {}",
                     self.get_size(*element_type, None),
-                    elem_align - 1,
-                    elem_align - 1,
+                    elem_align,
+                    elem_align,
+                    elem_align,
                     array_size
                 )
             }
@@ -2080,7 +2073,10 @@ namespace cg = cooperative_groups;
                         if acc == "0" {
                             size
                         } else {
-                            format!("(({} + {}) & ~{}) + {}", acc, align - 1, align - 1, size)
+                            format!(
+                                "({} + {} - 1) / {} * {} + {}",
+                                acc, align, align, align, size
+                            )
                         }
                     })
             }
@@ -2105,10 +2101,8 @@ namespace cg = cooperative_groups;
                     .max()
                     .unwrap_or(0);
                 format!(
-                    "(({} + {}) & ~{})",
-                    max_size,
-                    max_alignment - 1,
-                    max_alignment - 1,
+                    "({} + {} - 1) / {} * {}",
+                    max_size, max_alignment, max_alignment, max_alignment
                 )
             }
             _ => {