Incorrect SPIR-V shader passthrough

[HeroicKatora]

Passing SPIR-V shaders to naga and back alters their effects.

For example see this fragment shader used to paint a 2d normal distribution.
https://github.com/HeroicKatora/stealth-paint/blob/wgpu-0.10/src/shaders/distribution_normal2d.frag

We paint over a screen-filling quad (uv = [0; 1]×[0; 1]). With parameters

let u_fragmentParams_std140: &[f32; 13] = &[
    0.0, 0.0, 0.0, 0.0, 
    100.080055, 99.83977, 0.0, 0.0, 
    99.83977, 99.6012, 0.0, 0.0, 
    0.031466257
]

Expected:

Actual:

With parameters

let u_fragmentParams_std140: &[f32; 13] = &[
    0.0, 0.0, 0.0, 0.0, 
    5.0, 0.0, 0.0, 0.0, 
    0.0, 5.0, 0.0, 0.0, 
    1.5791368]

Expected:

Actual:

Additional information

Reproduce: Run the test suite here.

# Enables SPIR-V passthrough, works okay.
STEALTH_PAINT_PASSTHROUGH=1 STEALTH_PAINT_BLESS=1 cargo test
# Goes through Naga, fails
STEALTH_PAINT_BLESS=1 cargo test

There is a cryptic, reappearing validation error that seems to be unrelated.

[2021-09-20T07:51:41Z ERROR wgpu_core::validation] Unexpected varying type: Array { base: [1], size: Constant([11]), stride: 4 }

[kvark]

This is a great report!

[kvark]

I looked at this and found the problem in mat2 layout. In the original SPIR-V:

               OpMemberDecorate %FragmentColor 0 Offset 0
               OpMemberDecorate %FragmentColor 1 ColMajor
               OpMemberDecorate %FragmentColor 1 Offset 16
               OpMemberDecorate %FragmentColor 1 MatrixStride 16
               OpMemberDecorate %FragmentColor 2 Offset 48

In the produced SPIR-V:

               OpMemberDecorate %FragmentColor 0 Offset 0
               OpMemberDecorate %FragmentColor 1 Offset 16
               OpMemberDecorate %FragmentColor 1 ColMajor
               OpMemberDecorate %FragmentColor 1 MatrixStride 8
               OpMemberDecorate %FragmentColor 2 Offset 48

I'm surprised that the matrix stride on SPIR-V is 16. Here are the steps to proceed:

1. double-check if our logic wrt mat2 matches the WGSL spec right now (our validation and IR is based off WGSL largely)
2. decide if we want to allow weird strides in IR, or make SPV-in do a workaround to resolve this.

[kvark]

In the meantime, you can work around it by using vec4 instead of mat2.

[HeroicKatora]

That makes quite a lot of sense. How did you go about dumping the produced SPIR-V? This could be caused by being of the key differences in layout between std140 and std430. In the former, vec2 is padded to 16 bytes while in the latter those elements are contiguous in memory. Since mat2 is internally specified as being equivalent to vec2[2] we must pad the matrix rows in the former type. It might be that naga always assumes the unpadded layout.

[JCapucho]

This seems like a problem with spv-out we aren't passing the matrix stride and it's assuming the stride is the vector length, which is correct in std430 (also probably what wgsl uses but need to check) but since glsl defines that matrix stride must be a multiple of a vec4 alignment (16 machine units) in std140 it isn't working.

The long term solution is allowing to pass the matrix stride, the short term solution is algining with what wgsl and you potentially will need to use std430.

[HeroicKatora]

When I try to use std430 then I get the error that

"src/shaders/distribution_normal2d.frag:10: error: 'std430 requires the buffer storage qualifier' : required extension not requested: GL_EXT_scalar_block_layout\n"

(And I don't want to request any extensions above WebGPU). The workaround of vec4[2] works perfectly so I'm going to stick with that for now.

[JCapucho]

Yes glsl requires you to enable GL_EXT_scalar_block_layout to be able to use std430 in uniform blocks but as you're compiling to spirv this extension is always available so adding

#extension GL_EXT_scalar_block_layout : require

to the top of your shader should fix the error

[HeroicKatora]

Perfect, thank you.

[JCapucho]

@kvark I've checked and the wgsl spec defines that the matrices strides follow the std430 rules so the stride is that of the vector alignment (not rounded to 16 bytes). We could fix this in the spriv and glsl frontends, this would also extend to arrays since they have the same logic.

Another issue is the glsl backend because glsl requires the GL_EXT_scalar_block_layout extension for this strides but that extension requires GL_KHR_vulkan_glsl which is counter-productive since the whole idea is that we have no access to vulkan.

[kvark]

We have not need GL_EXT_scalar_block_layout at all.

Yes glsl requires you to enable GL_EXT_scalar_block_layout to be able to use std430 in uniform blocks

WGSL should work with std140 in uniform blocks. I think mat2 is the only tricky thing to get this going.
Can we replace it with vec4 in the backend and re-compose mat2 whenever we process Expression::Load?

[JCapucho]

WGSL should work with std140 in uniform blocks. I think mat2 is the only tricky thing to get this going. Can we replace it with vec4 in the backend and re-compose mat2 whenever we process Expression::Load?

It's not that easy because matrices, structures and arrays are affected by this rule of rounding the alignment to a vec4, for matrices we could potentially use a vec4 and for arrays we would need to do a swizzle if the member is anything smaller than a vec3, for structs I don't know because imagine this case where Outer is used as a uniform block:

struct Inner {
  float a; // 0 offset, 4 alignment
}

struct Outer {
  vec2 b; // 0 offset, 8 alignment
  Inner c; // 16 offset, 16 alignment because of round up
}

how do you access a? By wgsl rules c is residing at offset 8 but the generated glsl would try to access offset 16.

[kvark]

I don't think we need to do anything fancy about arrays. WGSL already requires that for uniform storage class:

Array elements are aligned to 16 byte boundaries. That is, StrideOf(array<T,N>) = 16 × k’ for some positive integer k’.

So naga expects incoming SPIR-V or whatever to already specify the stride=16 for such arrays, and the validation is happy, and the backend can just omit the stride if it sees this is a uniform buffer.

For structs, a similar rule is there:

If a structure member itself has a structure type S, then the number of bytes between the start of that member and the start of any following member must be at least roundUp(16, SizeOf(S)).

So this should mean WGSL uniform class is the same layout as std140.

Again, I think mat2 is the only exception.

[LPGhatguy]

I don't think this issue is limited to SPIR-V in or out. GLSL input also causes incorrect WGSL output!

Here's a simple pair of shaders to act as a litmus test. If both of these shaders result in the same WGSL, then this bug is still present:

#version 450

layout(std140, set = 0, binding = 0) readonly buffer INPUT {
    mat2 in_data;
};

layout(std140, set = 0, binding = 1) buffer OUTPUT {
    mat2 out_data;
};

void main() {
    out_data = in_data;
}

#version 450

layout(std430, set = 0, binding = 0) readonly buffer INPUT {
    mat2 in_data;
};

layout(std430, set = 0, binding = 1) buffer OUTPUT {
    mat2 out_data;
};

void main() {
    out_data = in_data;
}

These shaders differ only by their layout, which influences the layout of mat2. Currently, both shaders result in this WGSL:

[[block]]
struct INPUT {
    in_data: mat2x2<f32>;
};

[[block]]
struct OUTPUT {
    out_data: mat2x2<f32>;
};

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

fn main1() {
    let e4: mat2x2<f32> = global.in_data;
    global1.out_data = e4;
    return;
}

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

[kvark]

Maybe it's a separate issue that naga's GLSL frontend doesn't respect the layout qualifier? Related to #282, which got split and closed. @JCapucho looks like we are still missing this important bit - layout needs to change.

[JCapucho]

@kvark this is related to the lack of support for different matrix strides in the IR

[kvark]

Oh the mat2x2 thing? I think we can keep IR strict, but the GLSL and SPIR-V frontends would treat mat2x2 as vec4.

[teoxoy]

How should we deal with the other matrices of the form matCx2? (since they would also require a MatrixStride of 8 bytes (see resolution of gpuweb/gpuweb#1258))

[teoxoy]

I'm not sure that treating mat2x2 as vec4 is the right approach since that will essentially cause the buffer memory layout of the specific shading language (fontend) to no longer be reliable.

Imagine specifying std140, including a mat2 in that buffer, writing the data in the buffer according to the rules of std140 and then being surprised that the shader is seeing something else due to the mat2 now being read as a vec4.

[kvark]

I'm not sure that treating mat2x2 as vec4 is the right approach since that will essentially cause the buffer memory layout of the specific shading language (fontend) to no longer be reliable.

The frontend should make it legally and transparently. I.e. it wouldn't just change the type in IR, it would also change all of the operations on that type to correctly work with vec4.

This is doable, but also very low priority, since it's for GLSL and SPIR-V frontends only, a lot of work, and somewhat easy to work around for users.

[teoxoy]

What I was trying to point out is that for a buffer with the following data

offset	ty	mat2 (GLSL std140) interpretation	vec4 interpretation
0	f32	data0	data0
4	f32	data1	data1
8	f32	padding	data2
12	f32	padding	data3
16	f32	data2	? (determined by the next data type)
20	f32	data3	? (determined by the next data type)
24	f32	padding	? (determined by the next data type)
28	f32	padding	? (determined by the next data type)

the user will write a mat2 in the buffer as shown above. Reading the data as if it were a vec4 won't be right.

Unless we replace the mat2 (when in std140 or MatrixStride = 16) with a sequence of: vec2(data), vec2(padding), vec2(data), vec2(padding) and reconstruct it in the shader I don't see how this would work properly.

Having the user change the type manually works because they are expected to write the type into the buffer with the right layout but us changing it will only work if we read the data in the same way it was written.

Correct me if I'm wrong but I didn't see how the previous comments address this.

[teoxoy]

@kvark in light of your previous comment, I opened the following PRs/issues.

#1805 - [spv-in] error on an unsupported MatrixStride
#1806 - [glsl-in] error on a matCx2 used with the std140 layout
gfx-rs/wgpu#4375 - [spv-in] support any value of MatrixStride
gfx-rs/wgpu#4376 - [glsl-in] support matCx2 used with the std140 layout

I think we can close this issue and track the newly opened PRs/issues.

[kvark]

I see what you mean, thanks for clarification!