Eliminate empty struct on metal target

[kaizhangNV]

Close #6573.

We previously disabled the type legalization for ParameterBlock on Metal, but Metal doesn't allow empty struct in the argument buffer which is mapped from ParameterBlock, so we will need legalizeEmptyTypes on Metal target.

[csyonghe]

Why is this needed? Seems very hacky.

[kaizhangNV]

Because all the following decor check are based on C/C++.
For shader target, we always want to handle the type recursively. Otherwise, it will just return a LegalType::simple(type);

[kaizhangNV]

What really should be treated as non-simple type in this case is a type whose leaf type is a struct. However, we shouldn't repeat ourselves to do such check here, because there are already existing logic in legalizeTypeImpl() doing the same thing. We just want to make sure on Metal target, this will execute the logic in that function without early return.

[saipraveenb25]

Instead of adding a check inside the type legalization context, is is possible to create a different TypeLegalizationContext for the metal target, or not run legalizeEmptyTypes if the target is metal in slang-emit.cpp:1286?
I think we should try to keep all target-specific conditions at the top-level in linkAndOptimize so it's clear what's getting run/not run on different backends..

[kaizhangNV]

But I feel that the logic is almost 100% overlap with legalizeEmptyTypes. So there is no need to create different TypeLegalizationContext.

[kaizhangNV]

And, we need to run legalizeEmptyTypes if the target is metal, that is the point of this change actually.

[saipraveenb25]

Hmm..
I guess, since the legalization context does accept a TargetProgram, it's fine if target-specific stuff happens there.

[saipraveenb25]

Looks good!

[jkwak-work]

This PR seems to be causing CI failures.
https://github.com/shader-slang/slang/actions/runs/14092090746

The new test fails with the following error,

metal 32023.155: C:\Users\jkwak\AppData\Local\Temp\unkB06B.tmp.metal(30): error :  cannot initialize a member subobject of type 'int' with an lvalue of type 'device int *const constant'
    MyStruct_0 _S1 = { (*param_0).buffer_1, (*param_0).value_1 };
                       ^~~~~~~~~~~~~~~~~~~
1 error generated.

It looks like MyStruct_0 gets only one member variable whose type is int.
So it is an error to assign two values on the construction.

The following shader is the generated metal shader,

#include <metal_stdlib>
#include <metal_math>
#include <metal_texture>
using namespace metal;
void emptyStruct_set_0(int device* buffer_0, int value_0)
{
    *(buffer_0+int(0)) = value_0;
    return;
}

struct MyStruct_0
{
    int value_1;
};

void MyStruct_set_0(MyStruct_0 this_0, int device* this_buffer_0)
{
    emptyStruct_set_0(this_buffer_0, this_0.value_1);
    return;
}

struct MyStruct_1
{
    int device* buffer_1;
    int value_1;
};

[[kernel]] void computeMain(uint3 tid_0 [[thread_position_in_grid]], MyStruct_1 constant* param_0 [[buffer(0)]])
{
    MyStruct_0 _S1 = { (*param_0).buffer_1, (*param_0).value_1 };
    MyStruct_set_0(_S1, (*param_0).buffer_1);
    return;
}

I am wondering how this test passed CI, if it did.

[fairywreath]

The incorrect Metal code is caused by having RWStructuredBuffer<int> buffer; inside a struct. I have a temporary fix #6709 for the empty struct test that resolves this. Issue seems to be unrelated to the empty struct elimination, a separate fix to have RWStructuredBuffer inside a struct not emit incorrect metal code is required.