fixed mat2xc alignment in spv

CHANGELOG.md

@@ -63,6 +63,7 @@ Bottom level categories:
 #### Vulkan
 
 - Add support for RawWindowHandle::Drm on unix. By @rectalogic in [#9182](https://github.com/gfx-rs/wgpu/pull/9182).
+- Fixed alignment and MatrixStride for mat2x2 in SPIR-V uniform blocks. By @39ali [#9369](https://github.com/gfx-rs/wgpu/pull/9369).
 
 ### Changes
 

naga/src/back/spv/writer.rs

@@ -822,6 +822,9 @@ impl Writer {
         ir_module: &crate::Module,
         r#type: Handle<crate::Type>,
     ) -> Result<(), Error> {
+        if !self.std140_compat_uniform_types.contains_key(&r#type) {
+            return Ok(());
+        }
         // Check if we've already emitted this function.
         let wrapped = WrappedFunction::ConvertFromStd140CompatType { r#type };
         let function_id = match self.wrapped_functions.entry(wrapped) {
@@ -897,9 +900,7 @@ impl Writer {
                 self.write_wrapped_convert_from_std140_compat_type(ir_module, base)?;
 
                 let element_type_id = self.get_handle_type_id(base);
-                let std140_element_type_id = self.std140_compat_uniform_types[&base].type_id;
-                let element_conversion_function_id = self.wrapped_functions
-                    [&WrappedFunction::ConvertFromStd140CompatType { r#type: base }];
+                let std140_info = self.std140_compat_uniform_types.get(&base);
                 let mut element_ids = Vec::new();
                 let size = match size.resolve(ir_module.to_ctx())? {
                     crate::proc::IndexableLength::Known(size) => size,
@@ -911,20 +912,31 @@ impl Writer {
                 };
                 for i in 0..size {
                     let std140_element_id = self.id_gen.next();
+                    let std140_element_type_id =
+                        std140_info.map_or(element_type_id, |info| info.type_id);
                     block.body.push(Instruction::composite_extract(
                         std140_element_type_id,
                         std140_element_id,
                         param_id,
                         &[i],
                     ));
-                    let element_id = self.id_gen.next();
-                    block.body.push(Instruction::function_call(
-                        element_type_id,
-                        element_id,
-                        element_conversion_function_id,
-                        &[std140_element_id],
-                    ));
-                    element_ids.push(element_id);
+
+                    // Only call the conversion function if a compatibility mapping actually exists.
+                    let final_element_id = if std140_info.is_some() {
+                        let conversion_fn_id = self.wrapped_functions
+                            [&WrappedFunction::ConvertFromStd140CompatType { r#type: base }];
+                        let id = self.id_gen.next();
+                        block.body.push(Instruction::function_call(
+                            element_type_id,
+                            id,
+                            conversion_fn_id,
+                            &[std140_element_id],
+                        ));
+                        id
+                    } else {
+                        std140_element_type_id
+                    };
+                    element_ids.push(final_element_id);
                 }
                 let result_id = self.id_gen.next();
                 block.body.push(Instruction::composite_construct(
@@ -1000,7 +1012,6 @@ impl Writer {
                                     next_index += 1;
                                 }
                                 None => {
-                                    let member_id = self.id_gen.next();
                                     block.body.push(Instruction::composite_extract(
                                         member_type_id,
                                         member_id,

naga/tests/in/wgsl/mat2-uniform-alignment.toml

@@ -0,0 +1,2 @@
+
+targets = "SPIRV"

naga/tests/in/wgsl/mat2-uniform-alignment.wgsl

@@ -0,0 +1,11 @@
+struct S {
+  a : i32,
+  b : mat2x2<f32>,
+}
+
+@group(0) @binding(0) var<uniform> u : S;
+
+@compute @workgroup_size(1)
+fn main() {
+    let v = u;
+}

naga/tests/out/spv/wgsl-mat2-uniform-alignment.spvasm

@@ -0,0 +1,55 @@
+; SPIR-V
+; Version: 1.1
+; Generator: rspirv
+; Bound: 32
+OpCapability Shader
+%1 = OpExtInstImport "GLSL.std.450"
+OpMemoryModel Logical GLSL450
+OpEntryPoint GLCompute %22 "main"
+OpExecutionMode %22 LocalSize 1 1 1
+OpMemberDecorate %7 0 Offset 0
+OpMemberDecorate %7 1 Offset 8
+OpMemberDecorate %7 1 ColMajor
+OpMemberDecorate %7 1 MatrixStride 8
+OpMemberDecorate %8 0 Offset 0
+OpMemberDecorate %8 1 Offset 8
+OpMemberDecorate %8 2 Offset 16
+OpDecorate %9 DescriptorSet 0
+OpDecorate %9 Binding 0
+OpDecorate %10 Block
+OpMemberDecorate %10 0 Offset 0
+%2 = OpTypeVoid
+%3 = OpTypeInt 32 1
+%6 = OpTypeFloat 32
+%5 = OpTypeVector %6 2
+%4 = OpTypeMatrix %5 2
+%7 = OpTypeStruct %3 %4
+%8 = OpTypeStruct %3 %5 %5
+%10 = OpTypeStruct %8
+%11 = OpTypePointer Uniform %10
+%9 = OpVariable  %11  Uniform
+%13 = OpTypeFunction %7 %8
+%23 = OpTypeFunction %2
+%24 = OpTypePointer Uniform %8
+%26 = OpTypeInt 32 0
+%25 = OpConstant  %26  0
+%28 = OpTypePointer Uniform %7
+%12 = OpFunction  %7  None %13
+%14 = OpFunctionParameter  %8
+%15 = OpLabel
+%16 = OpCompositeExtract  %3  %14 0
+%18 = OpCompositeExtract  %5  %14 1
+%19 = OpCompositeExtract  %5  %14 2
+%17 = OpCompositeConstruct  %4  %18 %19
+%20 = OpCompositeConstruct  %7  %16 %17
+OpReturnValue %20
+OpFunctionEnd
+%22 = OpFunction  %2  None %23
+%21 = OpLabel
+%27 = OpAccessChain  %24  %9 %25
+OpBranch %29
+%29 = OpLabel
+%30 = OpLoad  %8  %27
+%31 = OpFunctionCall  %7  %12 %30
+OpReturn
+OpFunctionEnd