[MoltenVK][MSL] When `MVK_CONFIG_USE_METAL_ARGUMENT_BUFFERS` is `1`, IDs for descriptors in buffers are re-used, causing an error

[expenses]

I'm running a Vulkan program with MoltenVK. When I tried to use MVK_CONFIG_USE_METAL_ARGUMENT_BUFFERS=1 to increase performance, I ran into a shader compilation error.

The original SPIR-V shader is here: https://gist.github.com/expenses/c9ac3bad847cbcfeb709e9db0a5305a0

The translated metal shader is here (condensed as much as possible):

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

using namespace metal;

...

struct spvDescriptorSetBuffer0
{
    array<texture2d<float>, 127> m_95 [[id(0)]];
    sampler m_96 [[id(127)]];
    device _138* m_97 [[id(128)]];
    constant _139* m_98 [[id(129)]];
    sampler m_99 [[id(130)]];
};

struct spvDescriptorSetBuffer2
{
    constant uint* spvBufferSizeConstants [[id(0)]];
    device _140* m_100 [[id(0)]];
};

struct spvDescriptorSetBuffer3
{
    texture2d<float> m_101 [[id(0)]];
};

struct fragment_transmission_out
{
    float4 m_8 [[color(0)]];
};

struct fragment_transmission_in
{
    float3 m_3 [[user(locn0)]];
    float3 m_4 [[user(locn1)]];
    float2 m_5 [[user(locn2)]];
    uint m_6 [[user(locn3)]];
    float m_7 [[user(locn4), flat]];
};

fragment fragment_transmission_out fragment_transmission(fragment_transmission_in in [[stage_in]], constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], constant spvDescriptorSetBuffer2& spvDescriptorSet2 [[buffer(2)]], constant spvDescriptorSetBuffer3& spvDescriptorSet3 [[buffer(3)]], constant uint* spvBufferSizeConstants [[buffer(6)]], constant _137& _94 [[buffer(4)]])
{
    ...
}

The error I get is:

[mvk-error] VK_ERROR_INITIALIZATION_FAILED: Shader library compile failed (Error code 3):
program_source:129:201: error: cannot assign resource locations to 'spvDescriptorSet2'
fragment fragment_transmission_out fragment_transmission(fragment_transmission_in in [[stage_in]], constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], constant spvDescriptorSetBuffer2& spvDescriptorSet2 [[buffer(2)]], constant spvDescriptorSetBuffer3& spvDescriptorSet3 [[buffer(3)]], constant uint* spvBufferSizeConstants [[buffer(6)]], constant _137& _94 [[buffer(4)]])
                                                                                                                                                                                                        ^
program_source:107:18: note: attribute 'id' set location to 0, but minimum is 1
    device _140* m_100 [[id(0)]];
                 ^       ~~~~~

I don't know the exact steps that MoltenVK goes through to compile shaders but hopefully this is reproducable.

In case it helps, the full log of the program is here: https://gist.github.com/expenses/7a8e11b94b252db0607288d823a451fe

[expenses]

My guess is that when these spvDescriptorSetBuffer structs are created, the constant uint* spvBufferSizeConstants doesn't increase some id value when it really should.

[expenses]

Descriptor set 2 in this shader contains a single decriptor that is an unsized array (storage buffer in Vulkan) of structs - see https://github.com/expenses/transmission-renderer/blob/8323cdd5cc15ee762b135e5f248e40a6b7eeb218/shader/src/lib.rs#L48.

[adevaykin]

I think I have a triangle for this case.

GLSL code:

#define BINDING_PARAMBUFFER 0
#define BINDING_CHECKPOINT_DATA 1

struct CheckpointInterfaceDataParams
{
    uvec4 Data;
};

layout(std140, binding = BINDING_PARAMBUFFER) uniform CheckpointInterfaceData
{
    CheckpointInterfaceDataParams params;
};

struct CheckpointDataStruct
{
    uint lower;
    uint upper;
    uint phase;
    uint type;
};

layout(std430, binding=BINDING_CHECKPOINT_DATA) buffer CheckPoints
{
    uint count;
    uint padding[3];
    CheckpointDataStruct checkPoints[];
} Data;

layout(local_size_x = 1) in;
void main()
{
    if (gl_GlobalInvocationID.x == 0)
    {
        uint myIndex = atomicAdd(Data.count, 1);
        uint arrayIndex = myIndex % Data.checkPoints.length();
        Data.checkPoints[arrayIndex].lower = params.Data.x;
        Data.checkPoints[arrayIndex].upper = params.Data.y;
        Data.checkPoints[arrayIndex].phase = params.Data.z;
        Data.checkPoints[arrayIndex].type = params.Data.w;
    }
}

SPIR-V generation and crosscompiling to Metal:

glslc -std=430 -fshader-stage=compute shader.glsl
spirv-cross --msl a.spv --output shader.metal --msl-version 20100 --msl-argument-buffers

Resulting Metal shader:

#pragma clang diagnostic ignored "-Wunused-variable"

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

using namespace metal;

struct CheckpointDataStruct
{
    uint lower;
    uint upper;
    uint phase;
    uint type;
};

struct CheckPoints
{
    uint count;
    uint padding[3];
    CheckpointDataStruct checkPoints[1];
};

struct CheckpointInterfaceDataParams
{
    uint4 Data;
};

struct CheckpointInterfaceData
{
    CheckpointInterfaceDataParams params;
};

constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);

struct spvDescriptorSetBuffer0
{
    device CheckPoints* Data [[id(0)]];
    constant CheckpointInterfaceData* m_45 [[id(1)]];
    constant uint* spvBufferSizeConstants [[id(2)]];
};

kernel void main0(constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], constant uint* spvBufferSizeConstants [[buffer(25)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
{
    constant uint& spvDescriptorSet0_DataBufferSize = spvDescriptorSet0.spvBufferSizeConstants[0];
    if (gl_GlobalInvocationID.x == 0u)
    {
        uint _32 = atomic_fetch_add_explicit((device atomic_uint*)&(*spvDescriptorSet0.Data).count, 1u, memory_order_relaxed);
        uint myIndex = _32;
        uint arrayIndex = myIndex % uint(int((spvDescriptorSet0_DataBufferSize - 16) / 16));
        (*spvDescriptorSet0.Data).checkPoints[arrayIndex].lower = (*spvDescriptorSet0.m_45).params.Data.x;
        (*spvDescriptorSet0.Data).checkPoints[arrayIndex].upper = (*spvDescriptorSet0.m_45).params.Data.y;
        (*spvDescriptorSet0.Data).checkPoints[arrayIndex].phase = (*spvDescriptorSet0.m_45).params.Data.z;
        (*spvDescriptorSet0.Data).checkPoints[arrayIndex].type = (*spvDescriptorSet0.m_45).params.Data.w;
    }
}

See how device CheckPoints* Data [[id(0)]]; has id(0)