Array Stride Mismatch in vec3-padded Struct

[froody]

Description
Running a compute shader with an input of an array of struct of 20 bytes will only see the first 20/32 entries.

Repro steps

1. Clone https://github.com/froody/wgpu/tree/compute-bug
2. cd examples
3. cargo run hello_compute

Expected vs observed behavior
Observed:

result: 189
result with vertex 20: [0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 0, 0, 0, 0, 0, 0, 0, 0]
result with vertex 32: [0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19]

Expected:

result: 189
result with vertex 20: [0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19]
result with vertex 32: [0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19]

Extra materials
Screenshots to help explain your problem.
Validation logs can be attached in case there are warnings and errors.
Zip-compressed API traces and GPU captures can also land here.

Platform
MacOS 14.3 (23D56), MacBookPro18,2

[froody]

https://gist.github.com/froody/f1d4ec656a2110191ea1618187806ba1

[cwfitzgerald]

The issue here is as follows:

struct Inner { 
    vec: vec3<u32>,
    scalar1: u32,
    scalar2: u32
}

This struct is supposed to be size 32, alignment 16. However, when accessed through an array in metal, the stride shows itself to be 24. I suspect packed_uint3 doesn't have the alignment requirements?

struct Inner {
    metal::packed_uint3 vec;
    float scalar1;
    uint scalar2;
};

Unmark the test as failing on metal in #5264 to see the behavior.

[cwfitzgerald]

Yeah, https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf#page=31&zoom=auto,-462,785

[cwfitzgerald]

Probably just need to insert an alignas(16) when we're using packed_uint3