Translation of SPIR-V bindings into MSL is not correct

[Akaricchi]

I've addressed a special case of this in #23, but unfortunately this doesn't work in the general case and has caused a regression in TheSpydog/SDL_gpu_examples#46

The problem stems from the fact that MSL doesn't support descriptor sets. By default, SPIRV-Cross ignores the SPIR-V binding decoration. This makes it generate non-overlapping resource indices, but they are effectively unusable without some kind of reflection. With the change in #23, binding is preserved, but set is still ignored. This works for some simple cases, but breaks when there are multiple kinds of resources involved.

Example from the affected SpriteBatch compute shader:

layout (std430, set = 0, binding = 0) readonly buffer inBuffer
{
    SpriteComputeData computeData[];
};
layout (std430, set = 1, binding = 0) writeonly buffer outBuffer
{
    SpriteVertex vertexData[];
};

After #23 both get buffer index 0, causing a conflict:

2024-10-16 09:51:18.624 SDL_gpu_examples[34993:1465642] ERROR: Creating MTLLibrary failed: Error Domain=MTLLibraryErrorDomain Code=3 "program_source:41:79: error: cannot reserve 'buffer' resource location at index 0
kernel void main0(const device inBuffer& _27 [[buffer(0)]], device outBuffer& _114 [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
                                                                              ^

From SDL_GPU docs on resource layout for compute shaders:

For SPIR-V shaders, use the following resource sets:
- 0: Sampled textures, followed by read-only storage textures, followed by
  read-only storage buffers
- 1: Write-only storage textures, followed by write-only storage buffers
- 2: Uniform buffers

For MSL/metallib, use the following order:

- [[buffer]]: Uniform buffers, followed by write-only storage buffers,
  followed by write-only storage buffers
- [[texture]]: Sampled textures, followed by read-only storage textures,
  followed by write-only storage textures

(There seems to be a typo since it says write-ony storage twice; I'll assume it actually meant read-only followed by write-only)

We have to do some transformations to reconcile these differences:

• Subtract the number of sampled textures and read-only storage textures from the read-only storage buffer indices. Those live in the same descriptor set in SPIR-V and come before read-only storage buffers.
• Subtract the number of write-only storage textures from the write-only storage buffer indices. Same story.
• Add the number of read-only storage buffers to the write-only storage buffer indices, assuming ro comes before rw. Those share the same namespace in MSL.
• Add the number of uniform buffers to both read-only and write-only storage buffer indices. Those also live in the same namespace in MSL, and come first.
• Add the number of sampled textures and read-only storage textures to write-only storage texture indices. Again, SPIRV puts the write-only stuff into a separate descriptor set, while MSL has only two namespaces: buffers and textures.

A similar kind of dance has to be performed for vertex/fragment shaders as well, which have slightly different layout requirements.

It should be possible to do this by patching the SPIR-V before sending it to SPIRV-Cross, but it would be much easier (and beneficial to other projects) to implement these transformations in SPIRV-Cross itself. Basically we'd need to add an index offset option for each kind of resource.

An alternative way to solve this is to rework the Vulkan descriptor set layout to be more compatible with MSL at the SDL_GPU level. Unfortunately, that means breaking compatibility with existing SPIR-V shaders.

[thatcosmonaut]

I think for MSL we are intended to declare a remapping scheme, see CompilerMSL::add_msl_resource_binding in SPIRV-Cross.

[flibitijibibo]

In the meantime should we revert the MSL changes from #23? Not sure if that option is still needed when making the mapping scheme.

[Akaricchi]

It probably won't be needed, though by reverting it now you'd be trading one type of breakage for another (anything using multiple samplers or buffers is potentially affected). I've discovered the original problem because it broke some shaders in Taisei. I don't actually use shadercross for Taisei though, so it wouldn't affect me if you were to revert it.

[TheSpydog]

Since this change breaks the examples for the sake of a game that doesn't use shadercross, we've reverted this commit for now. We'll still definitely want to set up a remapping scheme as cosmonaut mentioned above.