Metal: support for structures in compute kernel arguments

[aroidzap]

Currently all structures in kernel arguments are flattened into arguments as separate structure members.

i.e. instead of:

struct MyStruct {
    RWStructuredBuffer<float> buffer;
    int value;
};

[[kernel]] void compute(uint3 globalIdx [[thread_position_in_grid]], MyStruct device* data);

this is generated:

[[kernel]] void compute(uint3 globalIdx [[thread_position_in_grid]], float device* buffer, int value);

Should be possible, see: https://developer.apple.com/documentation/realitykit/passing-structured-data-to-a-metal-compute-function

[csyonghe]

@aroidzap What you are asking is for slang to generate an argument buffer for MyStruct. This is done by doing:

void compute(ParameterBlock<MyStruct> data) {}

Otherwise, slang will hoist out all the resource types from the struct as top level parameters to fit the traditional binding model.

[aroidzap]

Thanks, I forgot about it, because it was not necessary when using [CudaKernel] or [DllExport].