Metal: Mesh Shaders

[Dynamitos]

Started working on the mesh shading stage, but this seems a lot more difficult than the task stage, because the IR is generated in a way where it references an output array and writes to it, but in Metal the output is set through a function call.

[csyonghe]

Would be good to have a test to exercise the new logic.

[Dynamitos]

This is just a basic idea for how the mesh output could be done, since the requirement to call a function instead of writing to an array is completely different. Primitives should work the same as vertices, indices could be easier since they have to be written to the array at once in HLSL, but haven't looked into it

[csyonghe]

A IRUse is a struct {IRInst user, IRInst usedVal;}

An IRInst is a op code + a tail-allocated array of IRUse. Each element IRUse in the array will have user be the inst itself, and usedVal be the inst referenced by the operand.

IRUse also has the functionality to allow you to iterate through all the other uses of "usedVal".

So if you have an IRUse* use, you can actually find the argument index of this operand by:

user = use->getUser();
operandId = use - user->getOperands();

The design here is the same as LLVM's IR def-use chain, and there may be documentation around that to explain this as well.

[csyonghe]

@Dynamitos Do you still plan to work on this patch? If so please reopen it so we don't lose track of it.

[Dynamitos]

I didnt get around to working on this for a while, and somehow messed up something with git, but now we should be back on track. I restarted from the master branch and integrated the changes again. I got a lot farther this time but this still needs to be done:

• Generate the MetalSetVertex, etc. IR conditionally in the core.meta.slang OR
• Replace the MeshOutputRef/MeshOutputSet with the metal versions in the legalize
• Remove the mesh array parameters after replacing the references with the mesh type
• Could probably use a bit of cleanup

[Dynamitos]

Also I noticed the [[position]] semantic isnt showing up, is this related to mesh shading or will this be fixed independently?

[csyonghe]

Also I noticed the [[position]] semantic isnt showing up, is this related to mesh shading or will this be fixed independently?

Do you have a test case that reproduces this issue? We may need to fix it separately if it can be reproduced with a vertex / fragment shader.

[Dynamitos]

It seems to work with the vertex shaders, but the simple-task.slang mesh stage test doesnt generate a position semantic, but user semantics are generated.

Also Im having trouble with a target_switch for the OutputVertices struct, i want to do something like this:

[require(glsl_hlsl_metal_spirv, meshshading)]
struct OutputVertices
{
    __subscript(uint index) -> T
    {
        // TODO: Make sure this remains write only, we can't do this with just
        // a 'set' operation as it's legal to only write to part of the output
        // buffer, or part of the output buffer at a time.
        __target_switch
        {
        case metal:
        {
            [mutating]
            __intrinsic_op($(kIROp_MetalSetVertex))
            set;
        }
        default:
        {
            __intrinsic_op($(kIROp_MeshOutputRef))
            ref;
        }
        }
    }
};

Since in metal the entire output needs to be written at once, this would directly generate the correct code, however it doesnt like the target_switch, do i have to escape it in some way?

[csyonghe]

You cannot use __target_switch outside a function body.

It should be:

[require(glsl_hlsl_metal_spirv, meshshading)]
struct OutputVertices
{
    __subscript(uint index) -> T
    {
        // TODO: Make sure this remains write only, we can't do this with just
        // a 'set' operation as it's legal to only write to part of the output
        // buffer, or part of the output buffer at a time.

            [mutating]
            set
            {
                    // implement set.
             }

             [require(glsl_hlsll_spirv)]
             __intrinsic_op($(kIROp_MeshOutputRef))
            __ref { ... }
    }
};

[csyonghe]

We should investigate why we need kIROp_MeshOutputRef at all. It seems like we can just use a getter and setter for both hlsl and spirv, and we make the getter require hlsl/spirv, while making the setter require hlsl/spirv/metal.

[Dynamitos]

if im inside of the set function body, the __intrinsic_op doesnt work anymore, and __intrinsic_asm only evaluates the IROp enum value.

for the position semantic problem i found out that it works if i specify the vertex struct as a return type, but i dont know how to properly fix this.

[csyonghe]

If the return type is not a struct, you will need to wrap it into one.

For intrinsic_op, you can always define a helper function that has an intrinsic_op modifier, then call that helper function inside the body of the setter.

[Dynamitos]

the problem is that the mesh function always returns void, and the vertex type which should contain the [[position]] semantic is just a parameter which can normally not have the semantic, the easiest solution i can think of is to use check the vertex type explicitly in the mesh stage instead of the return type, but this probably wont work for primitive attributes

[Dynamitos]

I haven fixed the position issue yet, and im not sure how to do that in a sane manner. Mesh outputs and payload should be correct now, however the other backends break when extracting the MeshOutputSet to a different function. I have tried to fix it, but I dont really understand the problem. But it definitely needs to be fixed before this can be merged. Alternatively the MeshOutputSets could be replaced in the legalize step, but this messes up the order of the code, the vertices get set before the values get initialized

[csyonghe]

@Dynamitos I still don't understand the position problem you are running into, can you help me understanding the problem with a concrete shader example and use it to explain what gap we are having right now?

[Dynamitos]

Here is an example:

struct Vert
{
    float4 pos : SV_Position;
};

struct Prim
{
    uint prim : SV_PrimitiveID;
}

[outputtopology("triangle")]
[numthreads(1, 1, 1)]
void test(OutputVertices<Vert, 3> verts, OutputIndices<uint3, 1> ind, OutputPrimitives<Prim, 1> prims)
{
    SetMeshOutputCounts(1, 1);
    verts[0] = { float4(0, 0, 0, 1) };
    verts[1] = { float4(1, 0, 0, 1) };
    verts[2] = { float4(0, 1, 0, 1) };
    ind[0] = uint3(0, 1, 2);
    prims[0] = { 0 };
}

This snippet produces the following metal code:

struct Vert_0
{
    float4 pos_0;
};
struct Prim_0
{
    uint prim_0;
};

[[mesh]] void test(metal::mesh<Vert_0, Prim_0, 3U, 1U, metal::topology::triangle> _slang_mesh)
{
    _slang_mesh.set_primitive_count((1U));
    Vert_0 _S1 = { float4(0.0, 0.0, 0.0, 1.0) };
    _slang_mesh.set_vertex(0U,_S1);
    Vert_0 _S2 = { float4(1.0, 0.0, 0.0, 1.0) };
    _slang_mesh.set_vertex(1U,_S2);
    Vert_0 _S3 = { float4(0.0, 1.0, 0.0, 1.0) };
    _slang_mesh.set_vertex(2U,_S3);
    _slang_mesh.set_index(0U*3,(uint3(0U, 1U, 2U))[0]);
    _slang_mesh.set_index(0U*3,(uint3(0U, 1U, 2U))[1]);
    _slang_mesh.set_index(0U*3,(uint3(0U, 1U, 2U))[2]);
    ;
    Prim_0 _S4 = { 0U };
    _slang_mesh.set_primitive(0U,_S4);
    return;
}

which is completely fine except for the fact that the [[position]] semantic is missing from the Vert struct and the [[primitive_id]] semantic is missing from the Prim struct.

[Dynamitos]

I fixed the semantics somewhat, not the prettiest code but it works for [[position]] and [[primitive_id]], cant think of any other outputs for the mesh shader stage.

[Dynamitos]

I have noticed two more things with the metal backend:

• SV_PrimitiveID is converted to [[patch_id]] when used as fragment shader input
• texture writes need to be converted to 4-component vectors regardless of the actual texture components.

The first one i fixed by replacing the metalSystemValueType, but it is odd enough to think that maybe there was a reason for it resoving to [[patch_id]]? For the second one it should probably be done in the legalization by inserting a MakeVector operation into the argument of the write.

Should this be handled in a separate pull request or merged with this one?

[csyonghe]

What should SV_PrimitiveID be translated into? It is likely a mistake in the original implementation. In that case a separate PR for that fix will be better, since we can get that merged in very quickly.

[Dynamitos]

there is a separate [[primitive_id]] semantic, [[patch_id]] is only valid after a tesselation stage

[csyonghe]

Yeah, that is a mistake and we should fix it in a separate PR.