MSL: Shader input overrides are not applied for block, array, and struct types

[spnda]

I found this during investigation of Vulkan CTS failures related KhronosGroup/MoltenVK#2116. For the following fragment inputs shader input overrides are not applied at all. Setting breakpoints on all usages of inputs_by_location shows that none of those codepaths are taken.

struct _14
{
    vec4 _m0;
    ivec3 _m1;
};

layout(location = 0) in _2
{
    flat _14 _m0[3];
} _19;

This generates the following MSL, regardless whether --msl-shader-input 1 any32 4 (which should widen the first int3 to int4) was specified. Also, the flat specifier is not being propagated properly.

struct main0_in
{
    float4 m_19_m0_0_m0 [[user(locn0)]];
    int3 m_19_m0_0_m1 [[user(locn1)]];
    float4 m_19_m0_1_m0 [[user(locn2)]];
    int3 m_19_m0_1_m1 [[user(locn3)]];
    float4 m_19_m0_2_m0 [[user(locn4)]];
    int3 m_19_m0_2_m1 [[user(locn5)]];
};

Note that this is a more complicated example, as this uses a struct within an array within a block, but as soon as any fragment inputs are within a block the shader inputs are ignored and not applied. Same happens with structs and arrays.

in _2
{
    layout(location = 0) vec2 _m0;
    layout(location = 1) flat uvec2 _m1;
} _15;

Shader outputs for the vertex shader using the same setup but for stage outputs are generated correctly.

I've attached a sample fragment shader that uses the above input setup:
17782185548125675323.spv.zip
This is the simplest example using just a block as shown above which also fails:
215374844162041042.spv.zip
And the tessellation shader with the same setup, where the shader outputs work correctly, if useful for debugging:
15478294089967604931.spv.zip

[HansKristian-Work]

Got a fix for the Flat propagation.

For the composites with override, there are many bugs. One of them is that OpLoad is too naive and only looks at the raw OpVariable.

		auto &expr_type = expression_type(ptr);

		// If the expression has more vector components than the result type, insert
		// a swizzle. This shouldn't happen normally on valid SPIR-V, but it might
		// happen with e.g. the MSL backend replacing the type of an input variable.
		if (expr_type.vecsize > type.vecsize)
			expr = enclose_expression(expr + vector_swizzle(type.vecsize, 0));

Need to rewrite all this junk ...

[dboyan]

Got a fix for the rest of mentioned issue in #2450. As long as the interface block is flattened, we just need to align all corresponding fields in the interface block with the specified with and add appropriate swizzles in the fixup. That actually resulted in a neat solution which doesn't require tweaking OpLoad.

But there is one missing case, i.e., when the interface block is not flattened for the input of tessellation eval shader. And honestly speaking, I don't have a good idea how to make it work yet.

Here is an example:

#version 450
layout(vertices = 1) out;

struct _29
{
    vec2 _m0;
    vec2 _m1;    // of type vec4 in the corresponding output from tesc stage
};

layout(location = 0) in _29 _33[];
layout(location = 0) out float _57[1];

void main()
{
    // ...
    float _27 = float(abs(_33[0]._m1.x - (-4.0)) < 0.001);
    _57[gl_InvocationID] = _27;
}

It translates to the following MSL with argument --msl-multi-patch-workgroup, regardless of how override is specified:

#include <metal_stdlib>
#include <simd/simd.h>

using namespace metal;

struct _29
{
    float2 _m0;
    float2 _m1;
};

struct main0_out
{
    float m_57;
};

struct main0_in
{
    _29 m_33;
};

kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]],
    device main0_out* spvOut [[buffer(28)]],
    constant uint* spvIndirectParams [[buffer(29)]],
    device MTLQuadTessellationFactorsHalf* spvTessLevel [[buffer(26)]],
    device main0_in* spvIn [[buffer(22)]])
{
    device main0_out* gl_out = &spvOut[gl_GlobalInvocationID.x - gl_GlobalInvocationID.x % 1];
    // INCORRECT, because the size of main0_in does not match output of previous stage
    // when there is an override
    device main0_in* gl_in = &spvIn[min(gl_GlobalInvocationID.x / 1, spvIndirectParams[1] - 1) * spvIndirectParams[0]];
    uint gl_InvocationID = gl_GlobalInvocationID.x % 1;
    // ...
    float _27 = float(abs(gl_in[0].m_33._m1.x - (-4.0)) < 0.001);
    gl_out[gl_InvocationID].m_57 = _27;
}

The problem here is that, in this case, input _33 is not flattened in the interface block main0_in. So the former is added in the latter with the method add_plain_variable_to_interface_block, which (I think) does not process overrides for composites. I don't have a clear view of how to properly handle this case. I don't think directly modifying the types in the interface block (e.g. _29 in the example above) is a good idea. Not sure if appropriate flattening can work there (on the input side).