HLSL backend uses wrong offsets for mat3x3<f32>

[LPGhatguy]

I debugged this one on the wgpu matrix server.

Feeding this WGSL shader into Naga:

struct TestData {
    one: mat3x3<f32>;
};

[[block]]
struct INPUT {
    in_data: TestData;
};

[[block]]
struct OUTPUT {
    out_data: TestData;
};

[[group(0), binding(0)]]
var<storage> global: INPUT;
[[group(0), binding(1)]]
var<storage, read_write> global1: OUTPUT;

fn main1() {
    let e4: TestData = global.in_data;
    global1.out_data = e4;
    return;
}

[[stage(compute), workgroup_size(1, 1, 1)]]
fn main() {
    main1();
    return;
}

Results in this HLSL:

struct TestData {
    row_major float3x3 one;
};

struct INPUT {
    TestData in_data;
};

struct OUTPUT {
    TestData out_data;
};

ByteAddressBuffer global : register(t0);
RWByteAddressBuffer global1_ : register(u1);

void main1_()
{
    TestData e4_ = {float3x3(asfloat(global.Load3(0+0+0)), asfloat(global.Load3(0+0+16)), asfloat(global.Load3(0+0+32)))};
    {
        TestData _value2 = e4_;
        {
            float3x3 _value3 = _value2.one;
            global1_.Store3(0+0+0, asuint(_value3[0]));
            global1_.Store3(0+0+12, asuint(_value3[1]));
            global1_.Store3(0+0+24, asuint(_value3[2]));
        }
    }
    return;
}

[numthreads(1, 1, 1)]
void main()
{
    main1_();
    return;
}

The problem is on lines 23-35 of the snippet:

global1_.Store3(0+0+0, asuint(_value3[0]));
global1_.Store3(0+0+12, asuint(_value3[1]));
global1_.Store3(0+0+24, asuint(_value3[2]));

These offsets are not correct for mat3 values. Curiously, the loads are correct!

This results in matrix3 values used in storage buffers being mangled on DX12, but not in Vulkan. Piping some data through this shader in wgpu results in:

Input:  TestData { one: ColumnMatrix3 { x: Vector3 { x: 1.0, y: 2.0, z: 3.0 }, y: Vector3 { x: 4.0, y: 5.0, z: 6.0 }, z: Vector3 { x: 7.0, y: 8.0, z: 9.0 } } }
Output: TestData { one: ColumnMatrix3 { x: Vector3 { x: 1.0, y: 2.0, z: 3.0 }, y: Vector3 { x: 5.0, y: 6.0, z: 7.0 }, z: Vector3 { x: 9.0, y: 0.0, z: 0.0 } } }