Repository: godotengine/godot (language: cpp)

Issue:
Metal: Error compiling compute shader on iPhone A12
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:

```metal
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:

https://github.com/godotengine/godot/blob/2450dee1bc1495daa2b7bdf06c7b3ddfbf8007e1/servers/rendering/renderer_rd/storage_rd/particles_storage.h#L157

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:

https://github.com/godotengine/godot/blob/d9ef826c545e7d623279c312ca837b947a4989b3/servers/rendering/renderer_rd/storage_rd/particles_storage.cpp#L544

_Originally posted by @TCROC in https://github.com/godotengine/godot/issues/99820#issuecomment-2558697999_
            


Discussion / hints:

Fix patch:
diff --git a/drivers/metal/metal_objects.mm b/drivers/metal/metal_objects.mm
index b44c496e4f6f..04900a0e5162 100644
--- a/drivers/metal/metal_objects.mm
+++ b/drivers/metal/metal_objects.mm
@@ -56,6 +56,10 @@
 
 #import <os/signpost.h>
 
+// We have to undefine these macros because they are defined in NSObjCRuntime.h.
+#undef MIN
+#undef MAX
+
 void MDCommandBuffer::begin() {
 	DEV_ASSERT(commandBuffer == nil);
 	commandBuffer = queue.commandBufferWithUnretainedReferences;
@@ -764,6 +768,7 @@
 		[render.encoder setVertexBuffers:render.vertex_buffers.ptr()
 								 offsets:render.vertex_offsets.ptr()
 							   withRange:NSMakeRange(first, p_binding_count)];
+		render.dirty.clear_flag(RenderState::DIRTY_VERTEX);
 	} else {
 		render.dirty.set_flag(RenderState::DIRTY_VERTEX);
 	}
@@ -1082,7 +1087,7 @@
 
 	UniformSet const &set = p_shader->sets[index];
 
-	for (uint32_t i = 0; i < uniforms.size(); i++) {
+	for (uint32_t i = 0; i < MIN(uniforms.size(), set.uniforms.size()); i++) {
 		RDD::BoundUniform const &uniform = uniforms[i];
 		UniformInfo ui = set.uniforms[i];
 
diff --git a/servers/rendering/renderer_rd/storage_rd/particles_storage.cpp b/servers/rendering/renderer_rd/storage_rd/particles_storage.cpp
index 7bd2d75c74a3..2ed4efd1572e 100644
--- a/servers/rendering/renderer_rd/storage_rd/particles_storage.cpp
+++ b/servers/rendering/renderer_rd/storage_rd/particles_storage.cpp
@@ -541,13 +541,15 @@ void ParticlesStorage::_particles_allocate_emission_buffer(Particles *particles)
 
 void ParticlesStorage::_particles_ensure_unused_emission_buffer(Particles *particles) {
 	if (particles->unused_emission_storage_buffer.is_null()) {
-		particles->unused_emission_storage_buffer = RD::get_singleton()->storage_buffer_create(sizeof(uint32_t) * 4);
+		// For rendering devices that do not support empty arrays (like C++),
+		// we need to size the buffer with at least 1 element.
+		particles->unused_emission_storage_buffer = RD::get_singleton()->storage_buffer_create(sizeof(ParticleEmissionBuffer));
 	}
 }
 
 void ParticlesStorage::_particles_ensure_unused_trail_buffer(Particles *particles) {
 	if (particles->unused_trail_storage_buffer.is_null()) {
-		particles->unused_trail_storage_buffer = RD::get_singleton()->storage_buffer_create(sizeof(uint32_t) * 4);
+		particles->unused_trail_storage_buffer = RD::get_singleton()->storage_buffer_create(16 * sizeof(float)); // Size of mat4.
 	}
 }
 


Test patch: