[glsl-out] Mat2xN is Misalligned in Uniform Buffers

[cwfitzgerald]

Found as part of gfx-rs/wgpu#3125

Mat2xN, in uniform buffers, are represented by glsl-out as literal matrices. But because of the spooky alignment at a distance that GLSL does, this actually ends up being the incorrect alignment. Just like HLSL, glsl-out needs to split the matrices into N different vec2s to get the correct alignment.

Test case:

struct InputStruct {
    member_0: mat2x2<f32>,
}

// The following removed by code in push contant case
@group(0) @binding(0)
var<uniform> input: InputStruct;

@group(0) @binding(1)
var<storage, read_write> output: array<f32>;

@compute @workgroup_size(1)
fn cs_main() {
    let loaded = input;
    var i = 0u;
    output[i] = input.member_0[0].x;
    i += 1u;
    output[i] = input.member_0[0].y;
    i += 1u;
    output[i] = input.member_0[1].x;
    i += 1u;
    output[i] = input.member_0[1].y;
    i += 1u;
}

Output GLSL from naga test.wgsl test.comp --entry-point cs_main:

#version 310 es

precision highp float;
precision highp int;

layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;

struct InputStruct {
    mat2x2 member_0_;
};
uniform InputStruct_block_0Compute { InputStruct _group_0_binding_0_cs; };

layout(std430) buffer type_2_block_1Compute { float _group_0_binding_1_cs[]; };


void main() {
    uint i = 0u;
    InputStruct loaded = _group_0_binding_0_cs;
    uint _e5 = i;
    float _e11 = _group_0_binding_0_cs.member_0_[0][0];
    _group_0_binding_1_cs[_e5] = _e11;
    uint _e12 = i;
    i = (_e12 + 1u);
    uint _e15 = i;
    float _e21 = _group_0_binding_0_cs.member_0_[0][1];
    _group_0_binding_1_cs[_e15] = _e21;
    uint _e22 = i;
    i = (_e22 + 1u);
    uint _e25 = i;
    float _e31 = _group_0_binding_0_cs.member_0_[1][0];
    _group_0_binding_1_cs[_e25] = _e31;
    uint _e32 = i;
    i = (_e32 + 1u);
    uint _e35 = i;
    float _e41 = _group_0_binding_0_cs.member_0_[1][1];
    _group_0_binding_1_cs[_e35] = _e41;
    uint _e42 = i;
    i = (_e42 + 1u);
    return;
}

[grovesNL]

This is probably the cause of gfx-rs/wgpu#3028

[teoxoy]

Duplicate of gfx-rs/wgpu#4371 but I guess we can keep both open for now.

[cwfitzgerald]

Closing in favor of gfx-rs/wgpu#4371 as I got both the title and body of the post wrong, though the test case still works