Transpilation of MSL from SPIRV compute shader with multiple input storage buffers generates unordered main0 signature

[MirkoCovizzi]

Input compute shader:

Archive.zip

Output msl:

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

using namespace metal;

struct instance_data
{
    float4x4 local_transform;
    uint vertex_count;
    uint vertex_offset;
};

struct instance_data_1
{
    float4x4 local_transform;
    uint vertex_count;
    uint vertex_offset;
    char _m0_final_padding[8];
};

struct in_buf_instance_data
{
    instance_data_1 data[1];
};

struct UBO
{
    float4x4 transform;
};

struct vertex_data
{
    float4 position;
    float2 tex_coord;
    float4 color;
};

struct out_buf
{
    vertex_data vertex_out[1];
};

struct in_buf_vertex_data
{
    vertex_data vertex_in[1];
};

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

kernel void main0(const device in_buf_instance_data& _26 [[buffer(0)]], constant UBO& ubo [[buffer(1)]], device out_buf& _75 [[buffer(2)]], const device in_buf_vertex_data& _81 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
{
    uint index = gl_GlobalInvocationID.x;
    instance_data current_data;
    current_data.local_transform = _26.data[index].local_transform;
    current_data.vertex_count = _26.data[index].vertex_count;
    current_data.vertex_offset = _26.data[index].vertex_offset;
    float4x4 model = ubo.transform * current_data.local_transform;
    uint vertex_offset = current_data.vertex_offset;
    for (uint i = 0u; i < current_data.vertex_count; i++)
    {
        uint global_vertex_index = vertex_offset + i;
        _75.vertex_out[global_vertex_index].position = model * _81.vertex_in[global_vertex_index].position;
        _75.vertex_out[global_vertex_index].tex_coord = _81.vertex_in[global_vertex_index].tex_coord;
        _75.vertex_out[global_vertex_index].color = _81.vertex_in[global_vertex_index].color;
    }
}

[TheSpydog]

Hmm, yeah, this is a mess. The original shader looks fine, but in the MSL we have, in order:

• readonly buffer
• uniform buffer
• read-write buffer
• readonly buffer

There was something we had to do to keep the order intact for the graphics pipeline in Metal, maybe something similar is needed for compute?

[thatcosmonaut]

I thought I had added a mapping system for compute as well but clearly it is not working as intended...