Implement manual particle emission and particle sub emitters.

[reduz]

Uses compute shader atomics to implement manual particle emission and sub emitters.

[bruvzg]

Shader compilation fails on MoltenVK:

ERROR:  - Message Id Number: 0 | Message Id Name:
	VK_ERROR_INITIALIZATION_FAILED: Shader library compile failed (Error code 3):
Compilation failed:

program_source:138:92: error: no viable overloaded operator[] for type 'volatile device metal::float4x4' (aka 'volatile device matrix<float, 4, 4>')
                    particles.data[particle].xform[3] = src_particles.data[src_index].xform[3];
                                                        ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~^~
program_source:147:92: error: no viable overloaded operator[] for type 'volatile device metal::float4x4' (aka 'volatile device matrix<float, 4, 4>')
                    particles.data[particle].xform[0] = src_particles.data[src_index].xform[0];
                                                        ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~^~
program_source:148:92: error: no viable overloaded operator[] for type 'volatile device metal::float4x4' (aka 'volatile device matrix<float, 4, 4>')
                    particles.data[particle].xform[1] = src_particles.data[src_index].xform[1];
                                                        ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~^~
program_source:149:92: error: no viable overloaded operator[] for type 'volatile device metal::float4x4' (aka 'volatile device matrix<float, 4, 4>')
                    particles.data[particle].xform[2] = src_particles.data[src_index].xform[2];
                                                        ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~^~

/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/3902/Libraries/lib/clang/902.14/include/metal/metal_matrix:277:35: note: candidate function not viable: 'this' argument ('volatile device metal::float4x4' (aka 'volatile device matrix<float, 4, 4>')) is in address space device, but parameter must be in address space 0
  METAL_FUNC thread vec<T, Rows> &operator[](int r) thread
                                  ^
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/3902/Libraries/lib/clang/902.14/include/metal/metal_matrix:281:35: note: candidate function not viable: 'this' argument has type 'volatile device metal::float4x4' (aka 'volatile device matrix<float, 4, 4>'), but method is not marked volatile
  METAL_FUNC device vec<T, Rows> &operator[](int r) device
                                  ^
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/3902/Libraries/lib/clang/902.14/include/metal/metal_matrix:285:40: note: candidate function not viable: 'this' argument ('volatile device metal::float4x4' (aka 'volatile device matrix<float, 4, 4>')) is in address space device, but parameter must be in address space threadgroup
  METAL_FUNC threadgroup vec<T, Rows> &operator[](int r) threadgroup
                                       ^
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/3902/Libraries/lib/clang/902.14/include/metal/metal_matrix:290:41: note: candidate function not viable: 'this' argument ('volatile device metal::float4x4' (aka 'volatile device matrix<float, 4, 4>')) is in address space device, but parameter must be in address space 0
  METAL_FUNC const thread vec<T, Rows> &operator[](int r) const thread
                                        ^
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/3902/Libraries/lib/clang/902.14/include/metal/metal_matrix:294:41: note: candidate function not viable: 'this' argument has type 'volatile device metal::float4x4' (aka 'volatile device matrix<float, 4, 4>'), but method is not marked volatile
  METAL_FUNC const device vec<T, Rows> &operator[](int r) const device
                                        ^
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/3902/Libraries/lib/clang/902.14/include/metal/metal_matrix:298:43: note: candidate function not viable: 'this' argument ('volatile device metal::float4x4' (aka 'volatile device matrix<float, 4, 4>')) is in address space device, but parameter must be in address space constant
  METAL_FUNC const constant vec<T, Rows> &operator[](int r) const constant
                                          ^
/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/3902/Libraries/lib/clang/902.14/include/metal/metal_matrix:302:46: note: candidate function not viable: 'this' argument ('volatile device metal::float4x4' (aka 'volatile device matrix<float, 4, 4>')) is in address space device, but parameter must be in address space threadgroup
  METAL_FUNC const threadgroup vec<T, Rows> &operator[](int r) const threadgroup
                                             ^

[reduz]

@bruvzg I thought empty arrays were allowed nowadays in C++, being 2020 but seems its not the case, so had to change the size to 1.

[bruvzg]

I thought empty arrays were allowed nowadays in C++, being 2020 but seems its not the case, so had to change the size to 1.

It's shader build error not engine build, zero size arrays do trigger warning but it's still building whiteout werror=yes (there's also invalid type uint pad that I have changed for my test build).

[akien-mga]

(there's also invalid type uint pad that I have changed for my test build)

Indeed:

D:\a\godot\godot\servers/rendering/rasterizer_rd/rasterizer_storage_rd.h(607): error C3646: 'pad': unknown override specifier
D:\a\godot\godot\servers/rendering/rasterizer_rd/rasterizer_storage_rd.h(607): error C4430: missing type specifier - int assumed. Note: C++ does not support default-int

[reduz]

ok lets hope it builds now

[bruvzg]

ok lets hope it builds now

It builds with werror=yes now, but the same shader error is still there.

[bruvzg]

Removing volatile and coherent qualifiers from here fixes shader compilation error, not sure what consequences this has. GPUParticles3D seems to be broken regardless of this change. It works with this change.

Suggested change

	layout(set = 1, binding = 2, std430) restrict volatile coherent buffer SourceEmission {
	layout(set = 1, binding = 2, std430) restrict buffer SourceEmission {