MSL: Argument buffer IDs are increased across buffers

[K0bin]

IDs of argument buffer members are local to each argument buffer. SPIR-V-Cross however keeps increasing them across argument buffers.

The GLSL snippet

struct AnotherStruct {
    uint _unused;
};
layout(set = 0, binding = 0, std430) readonly restrict buffer anotherBuffer {
  AnotherStruct arr[];
};

layout(set = 1, binding = 0) uniform texture2D textures[];

results in:

struct AnotherStruct
{
    uint _unused;
};

struct anotherBuffer
{
    AnotherStruct arr[1];
};

struct spvDescriptorSetBuffer1
{
    spvDescriptor<texture2d<float>> textures [[id(1)]][1] /* unsized array hack */;
};

There's no reason for the textures member to be [[id(1)]] instead of 0 and that causes problems when using it with an argument buffer that's manually encoded by writing the gpuResourceID to a buffer.

I'm using argument buffers tier 2 and all sets except the one that contains textures is set to be discrete (regular Metal bindings).

This attached shader reproduces it locally for me. (It's a bit different than the snippet above because I have a script that handles compiling GLSL all the way to a .mtllib file and that makes certain assumptions.) test.comp.glsl.zip

[HansKristian-Work]

How do you reproduce this? Seems to work for me.

./spirv-cross --msl /tmp/test.spv --msl-argument-buffers --msl-version 20000 --msl-device-argument-buffer 3 --msl-argument-buffer-tier 1

#pragma clang diagnostic ignored "-Wmissing-prototypes"

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

using namespace metal;

template<typename T>
struct spvDescriptor
{
    T value;
};

template<typename T>
struct spvDescriptorArray
{
    spvDescriptorArray(const device spvDescriptor<T>* ptr) : ptr(&ptr->value)
    {
    }
    const device T& operator [] (size_t i) const
    {
        return ptr[i];
    }
    const device T* ptr;
};

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

struct AnotherStruct
{
    uint _unused;
};

struct anotherBuffer
{
    AnotherStruct arr[1];
};

struct spvDescriptorSetBuffer0
{
    texture2d<float, access::write> image [[id(0)]];
    sampler linearSampler [[id(1)]];
};

struct spvDescriptorSetBuffer3
{
    spvDescriptor<texture2d<float>> textures [[id(0)]][1] /* unsized array hack */;
};

kernel void main0(constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], const device spvDescriptorSetBuffer3& spvDescriptorSet3 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
{
    spvDescriptorArray<texture2d<float>> textures {spvDescriptorSet3.textures};

    int2 texSize = int2(spvDescriptorSet0.image.get_width(), spvDescriptorSet0.image.get_height());
    bool _29 = gl_GlobalInvocationID.x >= uint(texSize.x);
    bool _40;
    if (!_29)
    {
        _40 = gl_GlobalInvocationID.y >= uint(texSize.y);
    }
    else
    {
        _40 = _29;
    }
    if (_40)
    {
        return;
    }
    float2 texCoord = float2((float(gl_GlobalInvocationID.x) + 0.5) / float(texSize.x), (float(gl_GlobalInvocationID.y) + 0.5) / float(texSize.y));
    int2 iTexCoord = int2(gl_GlobalInvocationID.xy);
    uint index = uint(iTexCoord.x);
    float3 color = textures[index].sample(spvDescriptorSet0.linearSampler, texCoord, level(0.0)).xyz;
    spvDescriptorSet0.image.write(float4(color, 1.0), uint2(iTexCoord));
}

[[id(0)]] is used here for the texture set.

[K0bin]

After discussing this on Discord, it turned out that the problem was me adding a resource binding for binding 0 of my bindless set using the SPIRV-Cross API. I thought this would set the Metal buffer binding that the argument buffer will use but instead it set the [[id(n)]] attribute of entry 0 of the argument buffer.