Metal: Error compiling compute shader on iPhone A12

[stuartcarnie]

When compiling the particles.glsl shader for an A12 or earlier device, it fails with the following error:

validateComputeFunctionArguments:1083: failed assertion `Compute Function(main0): argument src_particles[0] from buffer(4) with offset(0) and length(16) has space for 16 bytes, but argument has a length(128).'

Note

Disabling API validation allows the app to run.

The issue is that when using classic binding (i.e. not argument buffers), the ABI for the shader passes the structures as arguments to the entry point:

struct ParticleEmission
{
    float4x4 xform;
    packed_float3 velocity;
    uint flags;
    float4 color;
    float4 custom;
};

struct SourceEmission
{
    int particle_count;
    uint pad0;
    uint pad1;
    uint pad2;
    ParticleEmission data[1];
};

kernel void main0(device SourceEmission& src_particles [[buffer(5)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
{
// ...

Note that the src_particles argument is a struct SourceEmission. Note again that this struct has an embedded type struct ParticleEmission[1]. Metal doesn't support unsized arrays, but SPRIV-Cross works around it by passing an embedded struct of size 1. Interestingly, Metal is a variant of C++, and even Godot's definition has to use a sized array for the structure definition:

godot/servers/rendering/renderer_rd/storage_rd/particles_storage.h

Line 157 in 2450dee

Data data[1]; //its 2020 and empty arrays are still non standard in C++

The issue is that when Godot binds an empty emission buffer, it does not need to account for this, so it allocates a size 16:

godot/servers/rendering/renderer_rd/storage_rd/particles_storage.cpp

Line 544 in d9ef826

particles->unused_emission_storage_buffer = RD::get_singleton()->storage_buffer_create(sizeof(uint32_t) * 4);

Originally posted by @TCROC in #99820 (comment)