Fix type layout

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

hercules_cg/src/gpu.rs

@@ -2053,7 +2053,7 @@ namespace cg = cooperative_groups;
             Type::Array(element_type, extents) => {
                 assert!(num_fields.is_none());
                 let array_size = multiply_dcs(extents);
-                let elem_align = self.get_alignment(type_id);
+                let elem_align = self.get_alignment(*element_type);
                 format!(
                     "(({} + {} - 1) / {} * {}) * {}",
                     self.get_size(*element_type, None),