Fix type layout

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

hercules_cg/src/gpu.rs

@@ -2101,19 +2101,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 {