Metal: Mesh Shaders (#4280)

* Metal: mesh shading skeleton

* Metal: fixing mesh payload

* Metal: improving mesh shader indices output

* Metal: Implementing conditional mesh output set

* Metal: Trying to not break other backends

* Metal: trying to fix mesh output set

* Metal: Fixing MeshOutputSet usages

* Metal: Fixing vertex and primitive semantics

* Metal: Fixing code style

* Metal: Fixed hlsl indices set

* Fixed HLSL mesh output set disappearing and GLSL mesh output crashing

* Metal: Adjusting task test matching

---------

Co-authored-by: Yong He <yonghe@outlook.com>

1 files changed, 68 insertions, 2 deletions

diff --git a/source/slang/slang-emit-metal.cpp b/source/slang/slang-emit-metal.cpp
index 0b282c4d9..61e310401 100644
--- a/source/slang/slang-emit-metal.cpp
+++ b/source/slang/slang-emit-metal.cpp
@@ -492,6 +492,44 @@ bool MetalSourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inO
             m_writer->emit(")");
             return true;
         }
+        case kIROp_MetalSetVertex:
+        {
+            auto setVertex = as<IRMetalSetVertex>(inst);
+            m_writer->emit("_slang_mesh.set_vertex(");
+            emitOperand(setVertex->getIndex(), getInfo(EmitOp::General));
+            m_writer->emit(",");
+            emitOperand(setVertex->getElementValue(), getInfo(EmitOp::General));
+            m_writer->emit(")");
+            return true;
+        }
+        case kIROp_MetalSetPrimitive:
+        {
+            auto setPrimitive = as<IRMetalSetPrimitive>(inst);
+            m_writer->emit("_slang_mesh.set_primitive(");
+            emitOperand(setPrimitive->getIndex(), getInfo(EmitOp::General));
+            m_writer->emit(",");
+            emitOperand(setPrimitive->getElementValue(), getInfo(EmitOp::General));
+            m_writer->emit(")");
+            return true;
+        }
+        case kIROp_MetalSetIndices:
+        {
+            auto setIndices = as<IRMetalSetIndices>(inst);
+            const auto indices = as<IRVectorType>(setIndices->getElementValue()->getDataType());
+            UInt numIndices = as<IRIntLit>(indices->getElementCount())->getValue();
+            for(UInt i = 0; i < numIndices; ++i) {
+                m_writer->emit("_slang_mesh.set_index(");
+                emitOperand(setIndices->getIndex(), getInfo(EmitOp::General));
+                m_writer->emit("*");
+                m_writer->emitUInt64(numIndices);
+                m_writer->emit(",(");
+                emitOperand(setIndices->getElementValue(), getInfo(EmitOp::General));
+                m_writer->emit(")[");
+                m_writer->emitUInt64(i);
+                m_writer->emit("]);\n");
+            }
+            return true;
+        }
         default: break;
     }
     // Not handled
@@ -699,6 +737,10 @@ void MetalSourceEmitter::emitSimpleTypeImpl(IRType* type)
         case kIROp_ConstRefType:
         {
             auto ptrType = cast<IRPtrTypeBase>(type);
+            if(type->getOp() == kIROp_ConstRefType)
+            {
+                m_writer->emit("const ");
+            }
             emitType((IRType*)ptrType->getValueType());
             switch (ptrType->getAddressSpace())
             {
@@ -720,8 +762,7 @@ void MetalSourceEmitter::emitSimpleTypeImpl(IRType* type)
                 break;
             case AddressSpace::MetalObjectData:
                 m_writer->emit(" object_data");
-                // object data is passed by reference
-                m_writer->emit("&");
+                m_writer->emit("*");
                 break;
             }
             return;
@@ -785,6 +826,31 @@ void MetalSourceEmitter::emitSimpleTypeImpl(IRType* type)
         }
         return;
     }
+    else if (const auto meshType = as<IRMetalMeshType>(type))
+    {
+        m_writer->emit("metal::mesh<");
+        emitType(meshType->getVerticesType());
+        m_writer->emit(", ");
+        emitType(meshType->getPrimitivesType());
+        m_writer->emit(", ");
+        emitOperand(meshType->getNumVertices(), getInfo(EmitOp::General));
+        m_writer->emit(", ");
+        emitOperand(meshType->getNumPrimitives(), getInfo(EmitOp::General));
+        m_writer->emit(", metal::topology::");
+        switch(meshType->getTopology()->getValue()) {
+        case 1:
+            m_writer->emit("point");
+            break;
+        case 2:
+            m_writer->emit("line");
+            break;
+        case 3:
+            m_writer->emit("triangle");
+            break;
+        }
+        m_writer->emit(">");
+        return;
+    }
     else if(auto specializedType = as<IRSpecialize>(type))
     {
         // If a `specialize` instruction made it this far, then