HLSL to MSL : conversion between mismatching address spaces

[urbeller]

In a previously (working) translation, I had this MSL output (from hlsl compute shader):

struct type_RWByteAddressBuffer
{
    spvUnsafeArray<uint, 1> _m0;
};

struct type_OutputData
{
    uint2 outputSize;
    uint rowBytes;
    uint invertOrder;
};

constant uint2 _42 = {};

kernel void main0(device type_RWByteAddressBuffer& outputBuffer [[buffer(0)]], constant type_OutputData& OutputData [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
{

}

The same HLSL source with the recent cross-compiler led this following code:

struct type_RWByteAddressBuffer
{
    uint _m0[1];
};

struct type_OutputData
{
    uint2 outputSize;
    uint rowBytes;
    uint invertOrder;
};

kernel void main0(device void* spvBufferAliasSet0Binding0 [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
{
    device auto& outputBuffer = *(device type_RWByteAddressBuffer*)spvBufferAliasSet0Binding0;
    constant auto& OutputData = *(constant type_OutputData*)spvBufferAliasSet0Binding0;
...
}

First of all, I don't understand why buffer(1) suddenly disappeared ? is there an argument buffer trick that I am supposed to use ? Also, the new code generates the following error:

C-style cast from 'device void *' to 'const constant type_OutputData *' converts between mismatching address spaces

Update: this seems to occur because of register indexing clash. the device and constant buffers in hlsl were indexed as u0 and b0. Thank you,

[HansKristian-Work]

I'm confused. Did you resolve the issue yourself? If not, please upload reproducing SPIR-V.

[urbeller]

Sorry for the confusion. I haven't solved it but in deed I wasn't clear :-) . The following HLSL buffers:

RWStructuredBuffer<Vertex> vBuffer : register(u0);
cbuffer Params : register(b0)
{
    uint numVertices;
    float4 dims;
};

... is translated into Metal as:

kernel void main0(device void* spvBufferAliasSet0Binding0 [[buffer(0)]],...){

  device auto& vBuffer = *(device type_RWStructuredBuffer_Vertex*)spvBufferAliasSet0Binding0;
  constant auto& Params = *(constant type_Data*)spvBufferAliasSet0Binding0;
...
}

I realized that I am using the decorator to preserve the buffers indexing ! That could be the reason.

[HansKristian-Work]

You cannot alias UBO and SSBO like that. They have same binding. There is a way to do this now with indirect argument buffers, but not plain descriptors. MSL does not allow casting between address spaces like this.

[urbeller]

makes sens! thx for your help.