Fix type layout

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

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 {