format

* format

* Minor test fixes

* enable checking cpp format in ci

1 files changed, 430 insertions, 438 deletions

diff --git a/source/slang/slang-emit-metal.cpp b/source/slang/slang-emit-metal.cpp
index abd4d670a..ded0e05c2 100644
--- a/source/slang/slang-emit-metal.cpp
+++ b/source/slang/slang-emit-metal.cpp
@@ -2,14 +2,14 @@
 #include "slang-emit-metal.h"
 
 #include "../core/slang-writer.h"
-
-#include "slang-ir-util.h"
 #include "slang-emit-source-writer.h"
+#include "slang-ir-util.h"
 #include "slang-mangled-lexer.h"
 
 #include <assert.h>
 
-namespace Slang {
+namespace Slang
+{
 
 static const char* kMetalBuiltinPreludeMatrixCompMult = R"(
 template<typename T, int A, int B>
@@ -45,7 +45,10 @@ vec<T,A> _slang_vectorReshape(vec<T,N> v)
 }
 )";
 
-void MetalSourceEmitter::_emitHLSLDecorationSingleString(const char* name, IRFunc* entryPoint, IRStringLit* val)
+void MetalSourceEmitter::_emitHLSLDecorationSingleString(
+    const char* name,
+    IRFunc* entryPoint,
+    IRStringLit* val)
 {
     SLANG_UNUSED(entryPoint);
     assert(val);
@@ -57,7 +60,10 @@ void MetalSourceEmitter::_emitHLSLDecorationSingleString(const char* name, IRFun
     m_writer->emit("\")]]\n");
 }
 
-void MetalSourceEmitter::_emitHLSLDecorationSingleInt(const char* name, IRFunc* entryPoint, IRIntLit* val)
+void MetalSourceEmitter::_emitHLSLDecorationSingleInt(
+    const char* name,
+    IRFunc* entryPoint,
+    IRIntLit* val)
 {
     SLANG_UNUSED(entryPoint);
     SLANG_ASSERT(val);
@@ -71,9 +77,12 @@ void MetalSourceEmitter::_emitHLSLDecorationSingleInt(const char* name, IRFunc*
     m_writer->emit(")]]\n");
 }
 
-void MetalSourceEmitter::_emitHLSLParameterGroup(IRGlobalParam* varDecl, IRUniformParameterGroupType* type)
+void MetalSourceEmitter::_emitHLSLParameterGroup(
+    IRGlobalParam* varDecl,
+    IRUniformParameterGroupType* type)
 {
-    // Metal does not allow shader parameters declared as global variables, so we shouldn't see this.
+    // Metal does not allow shader parameters declared as global variables, so we shouldn't see
+    // this.
     SLANG_UNUSED(varDecl);
     SLANG_UNUSED(type);
     SLANG_ASSERT(!"Metal does not allow shader parameters declared as global variables.");
@@ -92,14 +101,12 @@ void MetalSourceEmitter::_emitHLSLTextureType(IRTextureTypeBase* texType)
 
     switch (texType->GetBaseShape())
     {
-        case SLANG_TEXTURE_1D:		m_writer->emit("1d");		break;
-        case SLANG_TEXTURE_2D:		m_writer->emit("2d");		break;
-        case SLANG_TEXTURE_3D:		m_writer->emit("3d");		break;
-        case SLANG_TEXTURE_CUBE:	m_writer->emit("cube");	    break;
-        case SLANG_TEXTURE_BUFFER:  m_writer->emit("_buffer");           break;
-        default:
-            SLANG_DIAGNOSE_UNEXPECTED(getSink(), SourceLoc(), "unhandled resource shape");
-            break;
+    case SLANG_TEXTURE_1D:     m_writer->emit("1d"); break;
+    case SLANG_TEXTURE_2D:     m_writer->emit("2d"); break;
+    case SLANG_TEXTURE_3D:     m_writer->emit("3d"); break;
+    case SLANG_TEXTURE_CUBE:   m_writer->emit("cube"); break;
+    case SLANG_TEXTURE_BUFFER: m_writer->emit("_buffer"); break;
+    default:                   SLANG_DIAGNOSE_UNEXPECTED(getSink(), SourceLoc(), "unhandled resource shape"); break;
     }
 
     if (texType->isMultisample())
@@ -113,20 +120,16 @@ void MetalSourceEmitter::_emitHLSLTextureType(IRTextureTypeBase* texType)
     m_writer->emit("<");
     emitType(getVectorElementType(texType->getElementType()));
     m_writer->emit(", ");
-    
+
     switch (texType->getAccess())
     {
-    case SLANG_RESOURCE_ACCESS_READ:
-        m_writer->emit("access::sample");
-        break;
+    case SLANG_RESOURCE_ACCESS_READ: m_writer->emit("access::sample"); break;
 
     case SLANG_RESOURCE_ACCESS_READ_WRITE:
     case SLANG_RESOURCE_ACCESS_APPEND:
     case SLANG_RESOURCE_ACCESS_CONSUME:
     case SLANG_RESOURCE_ACCESS_FEEDBACK:
-    case SLANG_RESOURCE_ACCESS_RASTER_ORDERED:
-        m_writer->emit("access::read_write");
-        break;
+    case SLANG_RESOURCE_ACCESS_RASTER_ORDERED: m_writer->emit("access::read_write"); break;
     default:
         SLANG_DIAGNOSE_UNEXPECTED(getSink(), SourceLoc(), "unhandled resource access mode");
         break;
@@ -149,7 +152,8 @@ void MetalSourceEmitter::emitFuncParamLayoutImpl(IRInst* param)
         switch (rr->getResourceKind())
         {
         case LayoutResourceKind::MetalTexture:
-            if (as<IRTextureTypeBase>(param->getDataType()) || as<IRTextureBufferType>(param->getDataType()))
+            if (as<IRTextureTypeBase>(param->getDataType()) ||
+                as<IRTextureBufferType>(param->getDataType()))
             {
                 m_writer->emit(" [[texture(");
                 m_writer->emit(rr->getOffset());
@@ -157,7 +161,8 @@ void MetalSourceEmitter::emitFuncParamLayoutImpl(IRInst* param)
             }
             break;
         case LayoutResourceKind::MetalBuffer:
-            if (as<IRPtrTypeBase>(param->getDataType()) || as<IRHLSLStructuredBufferTypeBase>(param->getDataType()) ||
+            if (as<IRPtrTypeBase>(param->getDataType()) ||
+                as<IRHLSLStructuredBufferTypeBase>(param->getDataType()) ||
                 as<IRByteAddressBufferTypeBase>(param->getDataType()) ||
                 as<IRUniformParameterGroupType>(param->getDataType()))
             {
@@ -174,12 +179,8 @@ void MetalSourceEmitter::emitFuncParamLayoutImpl(IRInst* param)
                 m_writer->emit(")]]");
             }
             break;
-        case LayoutResourceKind::VaryingInput:
-            m_writer->emit(" [[stage_in]]");
-            break;
-        case LayoutResourceKind::MetalPayload:
-            m_writer->emit(" [[payload]]");
-            break;
+        case LayoutResourceKind::VaryingInput: m_writer->emit(" [[stage_in]]"); break;
+        case LayoutResourceKind::MetalPayload: m_writer->emit(" [[payload]]"); break;
         }
     }
     if (!maybeEmitSystemSemantic(param))
@@ -189,34 +190,27 @@ void MetalSourceEmitter::emitFuncParamLayoutImpl(IRInst* param)
     }
 }
 
-void MetalSourceEmitter::emitParameterGroupImpl(IRGlobalParam* varDecl, IRUniformParameterGroupType* type)
+void MetalSourceEmitter::emitParameterGroupImpl(
+    IRGlobalParam* varDecl,
+    IRUniformParameterGroupType* type)
 {
     _emitHLSLParameterGroup(varDecl, type);
 }
 
-void MetalSourceEmitter::emitEntryPointAttributesImpl(IRFunc* irFunc, IREntryPointDecoration* entryPointDecor)
+void MetalSourceEmitter::emitEntryPointAttributesImpl(
+    IRFunc* irFunc,
+    IREntryPointDecoration* entryPointDecor)
 {
     auto stage = entryPointDecor->getProfile().getStage();
 
     switch (stage)
     {
-    case Stage::Fragment:
-        m_writer->emit("[[fragment]] ");
-        break;
-    case Stage::Vertex:
-        m_writer->emit("[[vertex]] ");
-        break;
-    case Stage::Compute:
-        m_writer->emit("[[kernel]] ");
-        break;
-    case Stage::Mesh:
-        m_writer->emit("[[mesh]] ");
-        break;
-    case Stage::Amplification:
-        m_writer->emit("[[object]] ");
-        break;
-    default:
-        SLANG_ABORT_COMPILATION("unsupported stage.");
+    case Stage::Fragment:      m_writer->emit("[[fragment]] "); break;
+    case Stage::Vertex:        m_writer->emit("[[vertex]] "); break;
+    case Stage::Compute:       m_writer->emit("[[kernel]] "); break;
+    case Stage::Mesh:          m_writer->emit("[[mesh]] "); break;
+    case Stage::Amplification: m_writer->emit("[[object]] "); break;
+    default:                   SLANG_ABORT_COMPILATION("unsupported stage.");
     }
 
     switch (stage)
@@ -229,8 +223,7 @@ void MetalSourceEmitter::emitEntryPointAttributesImpl(IRFunc* irFunc, IREntryPoi
             }
             break;
         }
-        default:
-            break;
+    default: break;
     }
 }
 
@@ -251,12 +244,8 @@ void MetalSourceEmitter::emitMemoryOrderOperand(IRInst* inst)
     auto memoryOrder = (IRMemoryOrder)getIntVal(inst);
     switch (memoryOrder)
     {
-    case kIRMemoryOrder_Relaxed:
-        m_writer->emit("memory_order_relaxed");
-        break;
-    default:
-        m_writer->emit("memory_order_seq_cst");
-        break;
+    case kIRMemoryOrder_Relaxed: m_writer->emit("memory_order_relaxed"); break;
+    default:                     m_writer->emit("memory_order_seq_cst"); break;
     }
 }
 
@@ -272,7 +261,10 @@ void MetalSourceEmitter::emitAtomicImageCoord(IRImageSubscript* inst)
     {
         if (as<IRVectorType>(textureType->getElementType()))
         {
-            getSink()->diagnose(inst, Diagnostics::unsupportedTargetIntrinsic, "atomic operation on non-scalar texture");
+            getSink()->diagnose(
+                inst,
+                Diagnostics::unsupportedTargetIntrinsic,
+                "atomic operation on non-scalar texture");
         }
     }
     bool isArray = getIntVal(resourceType->getIsArrayInst()) != 0;
@@ -282,17 +274,21 @@ void MetalSourceEmitter::emitAtomicImageCoord(IRImageSubscript* inst)
         if (auto coordType = as<IRVectorType>(inst->getCoord()->getDataType()))
         {
             m_writer->emit(".");
-            const char* elements[] = { "x", "y", "z", "w" };
+            const char* elements[] = {"x", "y", "z", "w"};
             for (IRIntegerValue i = 0; i < getIntVal(coordType->getElementCount()) - 1; i++)
                 m_writer->emit(elements[Math::Min(3, (int)i)]);
             m_writer->emit(", ");
             emitOperand(inst->getCoord(), getInfo(EmitOp::Postfix));
             m_writer->emit(".");
-            m_writer->emit(elements[Math::Min(3, (int)getIntVal(coordType->getElementCount()) - 1)]);
+            m_writer->emit(
+                elements[Math::Min(3, (int)getIntVal(coordType->getElementCount()) - 1)]);
         }
         else
         {
-            getSink()->diagnose(inst, Diagnostics::unsupportedTargetIntrinsic, "invalid image coordinate for atomic operation");
+            getSink()->diagnose(
+                inst,
+                Diagnostics::unsupportedTargetIntrinsic,
+                "invalid image coordinate for atomic operation");
         }
     }
     else
@@ -338,277 +334,278 @@ void MetalSourceEmitter::emitAtomicSrcOperand(bool isImage, IRInst* inst)
 bool MetalSourceEmitter::tryEmitInstStmtImpl(IRInst* inst)
 {
     auto emitAtomicOp = [&](const char* imageFunc, const char* bufferFunc)
-        {
-            emitInstResultDecl(inst);
-            bool isImageOp = false;
-            if (auto imageSubscript = isTextureAccess(inst))
-            {
-                emitOperand(imageSubscript->getImage(), getInfo(EmitOp::Postfix));
-                m_writer->emit(".");
-                m_writer->emit(imageFunc);
-                m_writer->emit("(");
-                emitAtomicImageCoord(imageSubscript);
-                isImageOp = true;
-            }
-            else
-            {
-                m_writer->emit(bufferFunc);
-                m_writer->emit("(");
-                emitAtomicDestOperand(inst->getOperand(0));
-            }
-            m_writer->emit(", ");
-            emitAtomicSrcOperand(isImageOp, inst->getOperand(1));
-            if (!isImageOp)
-            {
-                m_writer->emit(", ");
-                emitMemoryOrderOperand(inst->getOperand(inst->getOperandCount() - 1));
-            }
-            if (isImageOp)
-                m_writer->emit(").x;\n");
-            else
-                m_writer->emit(");\n");
-        };
-    auto diagnoseFloatAtommic = [&]()
-        {
-            getSink()->diagnose(inst, Diagnostics::unsupportedTargetIntrinsic, "floating point atomic operation");
-        };
-    switch (inst->getOp())
     {
-    case kIROp_discard:
-        m_writer->emit("discard_fragment();\n");
-        return true;
-    case kIROp_MetalAtomicCast:
-    {
-        auto oldValName = getName(inst);
-        auto op0 = inst->getOperand(0);
-
-        m_writer->emit("atomic_");
-        emitType(op0->getDataType());
-        m_writer->emit(" ");
-        m_writer->emit(oldValName);
-        m_writer->emit(" = ");
-
-        m_writer->emit("((atomic_");
-        emitType(op0->getDataType());
-        m_writer->emit(")(");
-        emitOperand(op0, getInfo(EmitOp::General));
-        m_writer->emit("));\n");
-        return true;
-    }
-    case kIROp_AtomicLoad:
-    {
-        if (isFloatingType(inst->getDataType()))
-            diagnoseFloatAtommic();
-
         emitInstResultDecl(inst);
         bool isImageOp = false;
         if (auto imageSubscript = isTextureAccess(inst))
         {
             emitOperand(imageSubscript->getImage(), getInfo(EmitOp::Postfix));
-            m_writer->emit(".atomic_load(");
+            m_writer->emit(".");
+            m_writer->emit(imageFunc);
+            m_writer->emit("(");
             emitAtomicImageCoord(imageSubscript);
             isImageOp = true;
         }
         else
         {
-            m_writer->emit("atomic_load_explicit(");
+            m_writer->emit(bufferFunc);
+            m_writer->emit("(");
             emitAtomicDestOperand(inst->getOperand(0));
         }
+        m_writer->emit(", ");
+        emitAtomicSrcOperand(isImageOp, inst->getOperand(1));
         if (!isImageOp)
         {
             m_writer->emit(", ");
-            emitMemoryOrderOperand(inst->getOperand(1));
+            emitMemoryOrderOperand(inst->getOperand(inst->getOperandCount() - 1));
         }
         if (isImageOp)
             m_writer->emit(").x;\n");
         else
             m_writer->emit(");\n");
-        return true;
-    }
-    case kIROp_AtomicStore:
+    };
+    auto diagnoseFloatAtommic = [&]()
     {
-        bool isImageOp = false;
-        if (auto imageSubscript = isTextureAccess(inst))
+        getSink()->diagnose(
+            inst,
+            Diagnostics::unsupportedTargetIntrinsic,
+            "floating point atomic operation");
+    };
+    switch (inst->getOp())
+    {
+    case kIROp_discard: m_writer->emit("discard_fragment();\n"); return true;
+    case kIROp_MetalAtomicCast:
         {
-            emitOperand(imageSubscript->getImage(), getInfo(EmitOp::Postfix));
-            m_writer->emit(".atomic_store(");
-            emitAtomicImageCoord(imageSubscript);
-            isImageOp = true;
+            auto oldValName = getName(inst);
+            auto op0 = inst->getOperand(0);
+
+            m_writer->emit("atomic_");
+            emitType(op0->getDataType());
+            m_writer->emit(" ");
+            m_writer->emit(oldValName);
+            m_writer->emit(" = ");
+
+            m_writer->emit("((atomic_");
+            emitType(op0->getDataType());
+            m_writer->emit(")(");
+            emitOperand(op0, getInfo(EmitOp::General));
+            m_writer->emit("));\n");
+            return true;
         }
-        else
+    case kIROp_AtomicLoad:
         {
-            m_writer->emit("atomic_store_explicit(");
-            emitAtomicDestOperand(inst->getOperand(0));
+            if (isFloatingType(inst->getDataType()))
+                diagnoseFloatAtommic();
+
+            emitInstResultDecl(inst);
+            bool isImageOp = false;
+            if (auto imageSubscript = isTextureAccess(inst))
+            {
+                emitOperand(imageSubscript->getImage(), getInfo(EmitOp::Postfix));
+                m_writer->emit(".atomic_load(");
+                emitAtomicImageCoord(imageSubscript);
+                isImageOp = true;
+            }
+            else
+            {
+                m_writer->emit("atomic_load_explicit(");
+                emitAtomicDestOperand(inst->getOperand(0));
+            }
+            if (!isImageOp)
+            {
+                m_writer->emit(", ");
+                emitMemoryOrderOperand(inst->getOperand(1));
+            }
+            if (isImageOp)
+                m_writer->emit(").x;\n");
+            else
+                m_writer->emit(");\n");
+            return true;
         }
-        m_writer->emit(", ");
-        emitAtomicSrcOperand(isImageOp, inst->getOperand(1));
-        if (!isImageOp)
+    case kIROp_AtomicStore:
         {
+            bool isImageOp = false;
+            if (auto imageSubscript = isTextureAccess(inst))
+            {
+                emitOperand(imageSubscript->getImage(), getInfo(EmitOp::Postfix));
+                m_writer->emit(".atomic_store(");
+                emitAtomicImageCoord(imageSubscript);
+                isImageOp = true;
+            }
+            else
+            {
+                m_writer->emit("atomic_store_explicit(");
+                emitAtomicDestOperand(inst->getOperand(0));
+            }
             m_writer->emit(", ");
-            emitMemoryOrderOperand(inst->getOperand(2));
+            emitAtomicSrcOperand(isImageOp, inst->getOperand(1));
+            if (!isImageOp)
+            {
+                m_writer->emit(", ");
+                emitMemoryOrderOperand(inst->getOperand(2));
+            }
+            m_writer->emit(");\n");
+            return true;
         }
-        m_writer->emit(");\n");
-        return true;
-    }
     case kIROp_AtomicExchange:
-    {
-        if (isFloatingType(inst->getDataType()))
-            diagnoseFloatAtommic();
+        {
+            if (isFloatingType(inst->getDataType()))
+                diagnoseFloatAtommic();
 
-        emitAtomicOp("atomic_exchange", "atomic_exchange_explicit");
-        return true;
-    }
+            emitAtomicOp("atomic_exchange", "atomic_exchange_explicit");
+            return true;
+        }
     case kIROp_AtomicCompareExchange:
-    {
-        if (isFloatingType(inst->getDataType()))
-            diagnoseFloatAtommic();
+        {
+            if (isFloatingType(inst->getDataType()))
+                diagnoseFloatAtommic();
 
-        bool isImageOp = false;
-        auto imageSubscript = isTextureAccess(inst);
-        isImageOp = (imageSubscript != nullptr);
+            bool isImageOp = false;
+            auto imageSubscript = isTextureAccess(inst);
+            isImageOp = (imageSubscript != nullptr);
 
-        emitType(inst->getDataType(), getName(inst));
-        m_writer->emit(";\n{\n");
-        if (isImageOp)
-            m_writer->emit("vec<");
-        emitType(inst->getDataType());
-        if (isImageOp)
-            m_writer->emit(", 4>");
-        m_writer->emit(" _metal_cas_comparand");
-        m_writer->emit(" = ");
-        emitOperand(inst->getOperand(1), getInfo(EmitOp::General));
-        m_writer->emit(";\n");
-        if (imageSubscript)
-        {
-            emitOperand(imageSubscript->getImage(), getInfo(EmitOp::Postfix));
-            m_writer->emit(".atomic_compare_exchange_weak(");
-            emitAtomicImageCoord(imageSubscript);
-        }
-        else
-        {
-            m_writer->emit("atomic_compare_exchange_weak_explicit(");
-            emitAtomicDestOperand(inst->getOperand(0));
-        }
-        m_writer->emit(", &_metal_cas_comparand, ");
-        emitAtomicSrcOperand(isImageOp, inst->getOperand(2));
-        if (!isImageOp)
-        {
-            m_writer->emit(", ");
-            emitMemoryOrderOperand(inst->getOperand(3));
-            m_writer->emit(", ");
-            emitMemoryOrderOperand(inst->getOperand(4));
+            emitType(inst->getDataType(), getName(inst));
+            m_writer->emit(";\n{\n");
+            if (isImageOp)
+                m_writer->emit("vec<");
+            emitType(inst->getDataType());
+            if (isImageOp)
+                m_writer->emit(", 4>");
+            m_writer->emit(" _metal_cas_comparand");
+            m_writer->emit(" = ");
+            emitOperand(inst->getOperand(1), getInfo(EmitOp::General));
+            m_writer->emit(";\n");
+            if (imageSubscript)
+            {
+                emitOperand(imageSubscript->getImage(), getInfo(EmitOp::Postfix));
+                m_writer->emit(".atomic_compare_exchange_weak(");
+                emitAtomicImageCoord(imageSubscript);
+            }
+            else
+            {
+                m_writer->emit("atomic_compare_exchange_weak_explicit(");
+                emitAtomicDestOperand(inst->getOperand(0));
+            }
+            m_writer->emit(", &_metal_cas_comparand, ");
+            emitAtomicSrcOperand(isImageOp, inst->getOperand(2));
+            if (!isImageOp)
+            {
+                m_writer->emit(", ");
+                emitMemoryOrderOperand(inst->getOperand(3));
+                m_writer->emit(", ");
+                emitMemoryOrderOperand(inst->getOperand(4));
+            }
+            m_writer->emit(");\n");
+            m_writer->emit(getName(inst));
+            m_writer->emit(" = _metal_cas_comparand");
+            if (isImageOp)
+                m_writer->emit(".x");
+            m_writer->emit(";\n}\n");
+            return true;
         }
-        m_writer->emit(");\n");
-        m_writer->emit(getName(inst));
-        m_writer->emit(" = _metal_cas_comparand");
-        if (isImageOp)
-            m_writer->emit(".x");
-        m_writer->emit(";\n}\n");
-        return true;
-    }
     case kIROp_AtomicAdd:
-    {
-        if (isFloatingType(inst->getDataType()))
-            diagnoseFloatAtommic();
+        {
+            if (isFloatingType(inst->getDataType()))
+                diagnoseFloatAtommic();
 
-        emitAtomicOp("atomic_fetch_add", "atomic_fetch_add_explicit");
-        return true;
-    }
+            emitAtomicOp("atomic_fetch_add", "atomic_fetch_add_explicit");
+            return true;
+        }
     case kIROp_AtomicSub:
-    {
-        if (isFloatingType(inst->getDataType()))
-            diagnoseFloatAtommic();
+        {
+            if (isFloatingType(inst->getDataType()))
+                diagnoseFloatAtommic();
 
-        emitAtomicOp("atomic_fetch_sub", "atomic_fetch_sub_explicit");
-        return true;
-    }
+            emitAtomicOp("atomic_fetch_sub", "atomic_fetch_sub_explicit");
+            return true;
+        }
     case kIROp_AtomicAnd:
-    {
-        emitAtomicOp("atomic_fetch_and", "atomic_fetch_and_explicit");
-        return true;
-    }
+        {
+            emitAtomicOp("atomic_fetch_and", "atomic_fetch_and_explicit");
+            return true;
+        }
     case kIROp_AtomicOr:
-    {
-        emitAtomicOp("atomic_fetch_or", "atomic_fetch_or_explicit");
-        return true;
-    }
-    case kIROp_AtomicXor:
-    {
-        emitAtomicOp("atomic_fetch_xor", "atomic_fetch_xor_explicit");
-        return true;
-    }
-    case kIROp_AtomicMin:
-    {
-        if (isFloatingType(inst->getDataType()))
-            diagnoseFloatAtommic();
-
-        emitAtomicOp("atomic_fetch_min", "atomic_fetch_min_explicit");
-        return true;
-    }
-    case kIROp_AtomicMax:
-    {
-        if (isFloatingType(inst->getDataType()))
-            diagnoseFloatAtommic();
-
-        emitAtomicOp("atomic_fetch_max", "atomic_fetch_max_explicit");
-        return true;
-    }
-    case kIROp_AtomicInc:
-    {
-        emitInstResultDecl(inst);
-        bool isImageOp = false;
-        if (auto imageSubscript = isTextureAccess(inst))
         {
-            emitOperand(imageSubscript->getImage(), getInfo(EmitOp::Postfix));
-            m_writer->emit(".atomic_fetch_add(");
-            emitAtomicImageCoord(imageSubscript);
-            isImageOp = true;
+            emitAtomicOp("atomic_fetch_or", "atomic_fetch_or_explicit");
+            return true;
         }
-        else
+    case kIROp_AtomicXor:
         {
-            m_writer->emit("atomic_fetch_add_explicit(");
-            emitAtomicDestOperand(inst->getOperand(0));
+            emitAtomicOp("atomic_fetch_xor", "atomic_fetch_xor_explicit");
+            return true;
         }
-        m_writer->emit(", 1");
-        if (!isImageOp)
+    case kIROp_AtomicMin:
         {
-            m_writer->emit(", ");
-            emitMemoryOrderOperand(inst->getOperand(1));
+            if (isFloatingType(inst->getDataType()))
+                diagnoseFloatAtommic();
+
+            emitAtomicOp("atomic_fetch_min", "atomic_fetch_min_explicit");
+            return true;
         }
-        if (isImageOp)
-            m_writer->emit(").x;\n");
-        else
-            m_writer->emit(");\n");
-        return true;
-    }
-    case kIROp_AtomicDec:
-    {
-        emitInstResultDecl(inst);
-        bool isImageOp = false;
-        if (auto imageSubscript = isTextureAccess(inst))
+    case kIROp_AtomicMax:
         {
-            emitOperand(imageSubscript->getImage(), getInfo(EmitOp::Postfix));
-            m_writer->emit(".atomic_fetch_sub(");
-            emitAtomicImageCoord(imageSubscript);
-            isImageOp = true;
+            if (isFloatingType(inst->getDataType()))
+                diagnoseFloatAtommic();
+
+            emitAtomicOp("atomic_fetch_max", "atomic_fetch_max_explicit");
+            return true;
         }
-        else
+    case kIROp_AtomicInc:
         {
-            m_writer->emit("atomic_fetch_sub_explicit(");
-            emitAtomicDestOperand(inst->getOperand(0));
+            emitInstResultDecl(inst);
+            bool isImageOp = false;
+            if (auto imageSubscript = isTextureAccess(inst))
+            {
+                emitOperand(imageSubscript->getImage(), getInfo(EmitOp::Postfix));
+                m_writer->emit(".atomic_fetch_add(");
+                emitAtomicImageCoord(imageSubscript);
+                isImageOp = true;
+            }
+            else
+            {
+                m_writer->emit("atomic_fetch_add_explicit(");
+                emitAtomicDestOperand(inst->getOperand(0));
+            }
+            m_writer->emit(", 1");
+            if (!isImageOp)
+            {
+                m_writer->emit(", ");
+                emitMemoryOrderOperand(inst->getOperand(1));
+            }
+            if (isImageOp)
+                m_writer->emit(").x;\n");
+            else
+                m_writer->emit(");\n");
+            return true;
         }
-        m_writer->emit(", 1");
-        if (!isImageOp)
+    case kIROp_AtomicDec:
         {
-            m_writer->emit(", ");
-            emitMemoryOrderOperand(inst->getOperand(1));
+            emitInstResultDecl(inst);
+            bool isImageOp = false;
+            if (auto imageSubscript = isTextureAccess(inst))
+            {
+                emitOperand(imageSubscript->getImage(), getInfo(EmitOp::Postfix));
+                m_writer->emit(".atomic_fetch_sub(");
+                emitAtomicImageCoord(imageSubscript);
+                isImageOp = true;
+            }
+            else
+            {
+                m_writer->emit("atomic_fetch_sub_explicit(");
+                emitAtomicDestOperand(inst->getOperand(0));
+            }
+            m_writer->emit(", 1");
+            if (!isImageOp)
+            {
+                m_writer->emit(", ");
+                emitMemoryOrderOperand(inst->getOperand(1));
+            }
+            if (isImageOp)
+                m_writer->emit(").x;\n");
+            else
+                m_writer->emit(");\n");
+            return true;
         }
-        if (isImageOp)
-            m_writer->emit(").x;\n");
-        else
-            m_writer->emit(");\n");
-        return true;
-    }
     }
     return false;
 }
@@ -617,9 +614,9 @@ bool MetalSourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inO
 {
     switch (inst->getOp())
     {
-        case kIROp_MakeVector:
-        case kIROp_MakeMatrix:
-        case kIROp_MakeVectorFromScalar:
+    case kIROp_MakeVector:
+    case kIROp_MakeMatrix:
+    case kIROp_MakeVectorFromScalar:
         {
             if (inst->getOperandCount() == 1)
             {
@@ -638,7 +635,7 @@ bool MetalSourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inO
             }
             break;
         }
-        case kIROp_MatrixReshape:
+    case kIROp_MatrixReshape:
         {
             ensurePrelude(kMetalBuiltinPreludeMatrixReshape);
             m_writer->emit("_slang_matrixReshape<");
@@ -651,7 +648,7 @@ bool MetalSourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inO
             m_writer->emit(")");
             return true;
         }
-        case kIROp_VectorReshape:
+    case kIROp_VectorReshape:
         {
             ensurePrelude(kMetalBuiltinPreludeVectorReshape);
             m_writer->emit("_slang_vectorReshape<");
@@ -662,15 +659,15 @@ bool MetalSourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inO
             m_writer->emit(")");
             return true;
         }
-        case kIROp_Mul:
+    case kIROp_Mul:
         {
             // Component-wise multiplication needs to be special cased,
             // because Metal uses infix `*` to express inner product
             // when working with matrices.
 
             // Are both operands matrices?
-            if (as<IRMatrixType>(inst->getOperand(0)->getDataType())
-                && as<IRMatrixType>(inst->getOperand(1)->getDataType()))
+            if (as<IRMatrixType>(inst->getOperand(0)->getDataType()) &&
+                as<IRMatrixType>(inst->getOperand(1)->getDataType()))
             {
                 ensurePrelude(kMetalBuiltinPreludeMatrixCompMult);
                 m_writer->emit("_slang_matrixCompMult(");
@@ -682,7 +679,7 @@ bool MetalSourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inO
             }
             break;
         }
-        case kIROp_FRem:
+    case kIROp_FRem:
         {
             m_writer->emit("fmod(");
             emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
@@ -691,21 +688,27 @@ bool MetalSourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inO
             m_writer->emit(")");
             return true;
         }
-        case kIROp_Select:
+    case kIROp_Select:
         {
             m_writer->emit("select(");
-            emitOperand(inst->getOperand(2), leftSide(getInfo(EmitOp::General), getInfo(EmitOp::General)));
+            emitOperand(
+                inst->getOperand(2),
+                leftSide(getInfo(EmitOp::General), getInfo(EmitOp::General)));
             m_writer->emit(", ");
-            emitOperand(inst->getOperand(1), leftSide(getInfo(EmitOp::General), getInfo(EmitOp::General)));
+            emitOperand(
+                inst->getOperand(1),
+                leftSide(getInfo(EmitOp::General), getInfo(EmitOp::General)));
             m_writer->emit(", ");
-            emitOperand(inst->getOperand(0), leftSide(getInfo(EmitOp::General), getInfo(EmitOp::General)));
+            emitOperand(
+                inst->getOperand(0),
+                leftSide(getInfo(EmitOp::General), getInfo(EmitOp::General)));
             m_writer->emit(")");
             return true;
         }
-        case kIROp_BitCast:
+    case kIROp_BitCast:
         {
             auto toType = inst->getDataType();
-            
+
             m_writer->emit("as_type<");
             emitType(toType);
             m_writer->emit(">(");
@@ -713,7 +716,7 @@ bool MetalSourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inO
             m_writer->emit(")");
             return true;
         }
-        case kIROp_StringLit:
+    case kIROp_StringLit:
         {
             const auto handler = StringEscapeUtil::getHandler(StringEscapeUtil::Style::Slang);
 
@@ -725,7 +728,7 @@ bool MetalSourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inO
 
             return true;
         }
-        case kIROp_ByteAddressBufferLoad:
+    case kIROp_ByteAddressBufferLoad:
         {
             // This only works for loads of 4-byte values.
             // Other element types should have been lowered by previous legalization passes.
@@ -741,7 +744,7 @@ bool MetalSourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inO
             m_writer->emit(")>>2])");
             return true;
         }
-        case kIROp_ByteAddressBufferStore:
+    case kIROp_ByteAddressBufferStore:
         {
             // This only works for loads of 4-byte values.
             // Other element types should have been lowered by previous legalization passes.
@@ -756,7 +759,7 @@ bool MetalSourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inO
             return true;
         }
         break;
-        case kIROp_RWStructuredBufferGetElementPtr:
+    case kIROp_RWStructuredBufferGetElementPtr:
         {
             EmitOpInfo outerPrec = inOuterPrec;
             bool needClose = false;
@@ -769,8 +772,8 @@ bool MetalSourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inO
             maybeCloseParens(needClose);
             return true;
         }
-        case kIROp_StructuredBufferLoad:
-        case kIROp_RWStructuredBufferLoad:
+    case kIROp_StructuredBufferLoad:
+    case kIROp_RWStructuredBufferLoad:
         {
             auto prec = getInfo(EmitOp::Postfix);
             emitOperand(inst->getOperand(0), leftSide(inOuterPrec, prec));
@@ -779,7 +782,7 @@ bool MetalSourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inO
             m_writer->emit("]");
             return true;
         }
-        case kIROp_RWStructuredBufferStore:
+    case kIROp_RWStructuredBufferStore:
         {
             auto prec = getInfo(EmitOp::Postfix);
             emitOperand(inst->getOperand(0), leftSide(inOuterPrec, prec));
@@ -789,18 +792,18 @@ bool MetalSourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inO
             emitOperand(inst->getOperand(2), getInfo(EmitOp::General));
             return true;
         }
-        case kIROp_ImageLoad:
+    case kIROp_ImageLoad:
         {
             auto imageOp = as<IRImageLoad>(inst);
             emitOperand(imageOp->getImage(), getInfo(EmitOp::General));
             m_writer->emit(".read(");
             emitOperand(imageOp->getCoord(), getInfo(EmitOp::General));
-            if(imageOp->hasAuxCoord1())
+            if (imageOp->hasAuxCoord1())
             {
                 m_writer->emit(",");
                 emitOperand(imageOp->getAuxCoord1(), getInfo(EmitOp::General));
             }
-            if(imageOp->hasAuxCoord2())
+            if (imageOp->hasAuxCoord2())
             {
                 m_writer->emit(",");
                 emitOperand(imageOp->getAuxCoord2(), getInfo(EmitOp::General));
@@ -808,7 +811,7 @@ bool MetalSourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inO
             m_writer->emit(")");
             return true;
         }
-        case kIROp_ImageStore:
+    case kIROp_ImageStore:
         {
             auto imageOp = as<IRImageStore>(inst);
             emitOperand(imageOp->getImage(), getInfo(EmitOp::General));
@@ -816,7 +819,7 @@ bool MetalSourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inO
             emitOperand(imageOp->getValue(), getInfo(EmitOp::General));
             m_writer->emit(",");
             emitOperand(imageOp->getCoord(), getInfo(EmitOp::General));
-            if(imageOp->hasAuxCoord1())
+            if (imageOp->hasAuxCoord1())
             {
                 m_writer->emit(",");
                 emitOperand(imageOp->getAuxCoord1(), getInfo(EmitOp::General));
@@ -824,7 +827,7 @@ bool MetalSourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inO
             m_writer->emit(")");
             return true;
         }
-        case kIROp_MetalSetVertex:
+    case kIROp_MetalSetVertex:
         {
             auto setVertex = as<IRMetalSetVertex>(inst);
             m_writer->emit("_slang_mesh.set_vertex(");
@@ -834,7 +837,7 @@ bool MetalSourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inO
             m_writer->emit(")");
             return true;
         }
-        case kIROp_MetalSetPrimitive:
+    case kIROp_MetalSetPrimitive:
         {
             auto setPrimitive = as<IRMetalSetPrimitive>(inst);
             m_writer->emit("_slang_mesh.set_primitive(");
@@ -844,12 +847,13 @@ bool MetalSourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inO
             m_writer->emit(")");
             return true;
         }
-        case kIROp_MetalSetIndices:
+    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) {
+            for (UInt i = 0; i < numIndices; ++i)
+            {
                 m_writer->emit("_slang_mesh.set_index(");
                 emitOperand(setIndices->getIndex(), getInfo(EmitOp::General));
                 m_writer->emit("*");
@@ -864,7 +868,7 @@ bool MetalSourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inO
             }
             return true;
         }
-        default: break;
+    default: break;
     }
     // Not handled
     return false;
@@ -909,17 +913,17 @@ static bool _canEmitExport(const Profile& profile)
     return (family == ProfileFamily::DX && version >= ProfileVersion::DX_6_1);
 }
 
-/* virtual */void MetalSourceEmitter::emitFuncDecorationsImpl(IRFunc* func)
+/* virtual */ void MetalSourceEmitter::emitFuncDecorationsImpl(IRFunc* func)
 {
     // Specially handle export, as we don't want to emit it multiple times
-    if (getTargetProgram()->getOptionSet().getBoolOption(CompilerOptionName::GenerateWholeProgram) && 
+    if (getTargetProgram()->getOptionSet().getBoolOption(
+            CompilerOptionName::GenerateWholeProgram) &&
         _canEmitExport(m_effectiveProfile))
     {
         for (auto decoration : func->getDecorations())
         {
             const auto op = decoration->getOp();
-            if (op == kIROp_PublicDecoration ||
-                op == kIROp_HLSLExportDecoration)
+            if (op == kIROp_PublicDecoration || op == kIROp_HLSLExportDecoration)
             {
                 m_writer->emit("export\n");
                 break;
@@ -953,33 +957,33 @@ void MetalSourceEmitter::emitSimpleValueImpl(IRInst* inst)
 {
     switch (inst->getOp())
     {
-        case kIROp_FloatLit:
+    case kIROp_FloatLit:
         {
             IRConstant* constantInst = static_cast<IRConstant*>(inst);
             IRConstant::FloatKind kind = constantInst->getFloatKind();
             switch (kind)
             {
-                case IRConstant::FloatKind::Nan:
+            case IRConstant::FloatKind::Nan:
                 {
                     m_writer->emit("(0.0 / 0.0)");
                     return;
                 }
-                case IRConstant::FloatKind::PositiveInfinity:
+            case IRConstant::FloatKind::PositiveInfinity:
                 {
                     m_writer->emit("(1.0 / 0.0)");
                     return;
                 }
-                case IRConstant::FloatKind::NegativeInfinity:
+            case IRConstant::FloatKind::NegativeInfinity:
                 {
                     m_writer->emit("(-1.0 / 0.0)");
                     return;
                 }
-                default: break;
+            default: break;
             }
             break;
         }
 
-        default: break;
+    default: break;
     }
 
     Super::emitSimpleValueImpl(inst);
@@ -994,48 +998,36 @@ void MetalSourceEmitter::emitSimpleTypeImpl(IRType* type)
 {
     switch (type->getOp())
     {
-        case kIROp_VoidType:
-        case kIROp_BoolType:
-        case kIROp_Int8Type:
-        case kIROp_IntType:
-        case kIROp_UInt8Type:
-        case kIROp_UIntType:
-        case kIROp_FloatType:
-        case kIROp_DoubleType:
-        case kIROp_HalfType:
+    case kIROp_VoidType:
+    case kIROp_BoolType:
+    case kIROp_Int8Type:
+    case kIROp_IntType:
+    case kIROp_UInt8Type:
+    case kIROp_UIntType:
+    case kIROp_FloatType:
+    case kIROp_DoubleType:
+    case kIROp_HalfType:
         {
             m_writer->emit(getDefaultBuiltinTypeName(type->getOp()));
             return;
         }
-        case kIROp_Int64Type:
-            m_writer->emit("long");
-            return;
-        case kIROp_UInt64Type:
-            m_writer->emit("ulong");
-            return;
-        case kIROp_Int16Type:
-            m_writer->emit("short");
-            return;
-        case kIROp_UInt16Type:
-            m_writer->emit("ushort");
-            return;
-        case kIROp_IntPtrType:
-            m_writer->emit("long");
-            return;
-        case kIROp_UIntPtrType:
-            m_writer->emit("ulong");
-            return;
-        case kIROp_StructType:
-            m_writer->emit(getName(type));
-            return;
-
-        case kIROp_VectorType:
+    case kIROp_Int64Type:   m_writer->emit("long"); return;
+    case kIROp_UInt64Type:  m_writer->emit("ulong"); return;
+    case kIROp_Int16Type:   m_writer->emit("short"); return;
+    case kIROp_UInt16Type:  m_writer->emit("ushort"); return;
+    case kIROp_IntPtrType:  m_writer->emit("long"); return;
+    case kIROp_UIntPtrType: m_writer->emit("ulong"); return;
+    case kIROp_StructType:  m_writer->emit(getName(type)); return;
+
+    case kIROp_VectorType:
         {
             auto vecType = (IRVectorType*)type;
-            emitVectorTypeNameImpl(vecType->getElementType(), getIntVal(vecType->getElementCount()));
+            emitVectorTypeNameImpl(
+                vecType->getElementType(),
+                getIntVal(vecType->getElementCount()));
             return;
         }
-        case kIROp_MatrixType:
+    case kIROp_MatrixType:
         {
             auto matType = (IRMatrixType*)type;
 
@@ -1046,36 +1038,36 @@ void MetalSourceEmitter::emitSimpleTypeImpl(IRType* type)
             emitVal(matType->getRowCount(), getInfo(EmitOp::General));
             m_writer->emit(",");
             emitVal(matType->getColumnCount(), getInfo(EmitOp::General));
-            m_writer->emit("> ");           
+            m_writer->emit("> ");
             return;
         }
-        case kIROp_SamplerStateType:
-        case kIROp_SamplerComparisonStateType:
+    case kIROp_SamplerStateType:
+    case kIROp_SamplerComparisonStateType:
         {
             m_writer->emit("sampler");
             return;
         }
-        case kIROp_NativeStringType:
-        case kIROp_StringType: 
+    case kIROp_NativeStringType:
+    case kIROp_StringType:
         {
-            m_writer->emit("int"); 
+            m_writer->emit("int");
             return;
         }
-        case kIROp_ParameterBlockType:
-        case kIROp_ConstantBufferType:
+    case kIROp_ParameterBlockType:
+    case kIROp_ConstantBufferType:
         {
             emitSimpleTypeImpl((IRType*)type->getOperand(0));
             m_writer->emit(" constant*");
             return;
         }
-        case kIROp_PtrType:
-        case kIROp_InOutType:
-        case kIROp_OutType:
-        case kIROp_RefType:
-        case kIROp_ConstRefType:
+    case kIROp_PtrType:
+    case kIROp_InOutType:
+    case kIROp_OutType:
+    case kIROp_RefType:
+    case kIROp_ConstRefType:
         {
             auto ptrType = cast<IRPtrTypeBase>(type);
-            if(type->getOp() == kIROp_ConstRefType)
+            if (type->getOp() == kIROp_ConstRefType)
             {
                 m_writer->emit("const ");
             }
@@ -1105,7 +1097,7 @@ void MetalSourceEmitter::emitSimpleTypeImpl(IRType* type)
             }
             return;
         }
-        case kIROp_ArrayType:
+    case kIROp_ArrayType:
         {
             m_writer->emit("array<");
             emitType((IRType*)type->getOperand(0));
@@ -1114,20 +1106,19 @@ void MetalSourceEmitter::emitSimpleTypeImpl(IRType* type)
             m_writer->emit(">");
             return;
         }
-        case kIROp_MetalMeshGridPropertiesType:
+    case kIROp_MetalMeshGridPropertiesType:
         {
             m_writer->emit("mesh_grid_properties ");
             return;
         }
-        case kIROp_AtomicType:
+    case kIROp_AtomicType:
         {
             m_writer->emit("atomic<");
             emitSimpleTypeImpl(cast<IRAtomicType>(type)->getElementType());
             m_writer->emit(">");
             return;
         }
-        default:
-            break;
+    default: break;
     }
 
     if (auto texType = as<IRTextureType>(type))
@@ -1157,17 +1148,15 @@ void MetalSourceEmitter::emitSimpleTypeImpl(IRType* type)
     {
         switch (type->getOp())
         {
-            case kIROp_HLSLByteAddressBufferType:
-            case kIROp_HLSLRWByteAddressBufferType:
-            case kIROp_HLSLRasterizerOrderedByteAddressBufferType:
-                m_writer->emit("uint32_t device*");
-                break;
-            case kIROp_RaytracingAccelerationStructureType:
-                m_writer->emit("acceleration_structure<instancing>");
-                break;
-            default:
-                SLANG_DIAGNOSE_UNEXPECTED(getSink(), SourceLoc(), "unhandled buffer type");
-                break;
+        case kIROp_HLSLByteAddressBufferType:
+        case kIROp_HLSLRWByteAddressBufferType:
+        case kIROp_HLSLRasterizerOrderedByteAddressBufferType:
+            m_writer->emit("uint32_t device*");
+            break;
+        case kIROp_RaytracingAccelerationStructureType:
+            m_writer->emit("acceleration_structure<instancing>");
+            break;
+        default: SLANG_DIAGNOSE_UNEXPECTED(getSink(), SourceLoc(), "unhandled buffer type"); break;
         }
         return;
     }
@@ -1182,31 +1171,27 @@ void MetalSourceEmitter::emitSimpleTypeImpl(IRType* type)
         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;
+        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))
+    else if (auto specializedType = as<IRSpecialize>(type))
     {
         // If a `specialize` instruction made it this far, then
         // it represents an intrinsic generic type.
         //
-        emitSimpleType((IRType*) getSpecializedValue(specializedType));
+        emitSimpleType((IRType*)getSpecializedValue(specializedType));
         m_writer->emit("<");
         UInt argCount = specializedType->getArgCount();
         for (UInt ii = 0; ii < argCount; ++ii)
         {
-            if (ii != 0) m_writer->emit(", ");
+            if (ii != 0)
+                m_writer->emit(", ");
             emitVal(specializedType->getArg(ii), getInfo(EmitOp::General));
         }
         m_writer->emit(" >");
@@ -1224,7 +1209,8 @@ void MetalSourceEmitter::emitSimpleTypeImpl(IRType* type)
             m_writer->emit("<");
             for (UInt ii = 0; ii < operandCount; ++ii)
             {
-                if (ii != 0) m_writer->emit(", ");
+                if (ii != 0)
+                    m_writer->emit(", ");
                 emitVal(type->getOperand(ii), getInfo(EmitOp::General));
             }
             m_writer->emit(" >");
@@ -1240,9 +1226,7 @@ void MetalSourceEmitter::_emitType(IRType* type, DeclaratorInfo* declarator)
         emitSimpleType(type);
         emitDeclarator(declarator);
         break;
-    default:
-        Super::_emitType(type, declarator);
-        break;
+    default: Super::_emitType(type, declarator); break;
     }
 }
 
@@ -1258,7 +1242,9 @@ bool MetalSourceEmitter::maybeEmitSystemSemantic(IRInst* inst)
     return false;
 }
 
-bool MetalSourceEmitter::_emitUserSemantic(UnownedStringSlice semanticName, IRIntegerValue semanticIndex)
+bool MetalSourceEmitter::_emitUserSemantic(
+    UnownedStringSlice semanticName,
+    IRIntegerValue semanticIndex)
 {
     if (!semanticName.startsWithCaseInsensitive(toSlice("SV_")))
     {
@@ -1281,7 +1267,7 @@ void MetalSourceEmitter::emitSemanticsImpl(IRInst* inst, bool allowOffsets)
     if (inst->getOp() == kIROp_StructKey)
     {
         // Only emit [[attribute(n)]] on struct keys.
-        
+
         if (maybeEmitSystemSemantic(inst))
             return;
 
@@ -1305,25 +1291,29 @@ void MetalSourceEmitter::emitSemanticsImpl(IRInst* inst, bool allowOffsets)
             }
             for (auto attr : varLayout->getAllAttrs())
             {
-               if (auto semanticAttr = as<IRSemanticAttr>(attr))
+                if (auto semanticAttr = as<IRSemanticAttr>(attr))
                 {
                     auto semanticName = String(semanticAttr->getName()).toUpper();
-                    hasSemantic = _emitUserSemantic(semanticAttr->getName(), semanticAttr->getIndex());
+                    hasSemantic =
+                        _emitUserSemantic(semanticAttr->getName(), semanticAttr->getIndex());
                 }
             }
-
         }
         if (!hasSemantic)
         {
             if (auto semanticDecor = inst->findDecoration<IRSemanticDecoration>())
             {
-                _emitUserSemantic(semanticDecor->getSemanticName(), semanticDecor->getSemanticIndex());
+                _emitUserSemantic(
+                    semanticDecor->getSemanticName(),
+                    semanticDecor->getSemanticIndex());
             }
         }
     }
 }
 
-void MetalSourceEmitter::_emitStageAccessSemantic(IRStageAccessDecoration* decoration, const char* name)
+void MetalSourceEmitter::_emitStageAccessSemantic(
+    IRStageAccessDecoration* decoration,
+    const char* name)
 {
     SLANG_UNUSED(decoration);
     SLANG_UNUSED(name);
@@ -1345,7 +1335,8 @@ void MetalSourceEmitter::emitPostDeclarationAttributesForType(IRInst* type)
             m_writer->emit(" [[raster_order_group(0)]]");
         }
     }
-    else if (as<IRHLSLRasterizerOrderedByteAddressBufferType>(type) ||
+    else if (
+        as<IRHLSLRasterizerOrderedByteAddressBufferType>(type) ||
         as<IRHLSLRasterizerOrderedStructuredBufferType>(type))
     {
         m_writer->emit(" [[raster_order_group(0)]]");
@@ -1356,17 +1347,24 @@ static UnownedStringSlice _getInterpolationModifierText(IRInterpolationMode mode
 {
     switch (mode)
     {
-        case IRInterpolationMode::PerVertex:
-        case IRInterpolationMode::NoInterpolation:      return UnownedStringSlice::fromLiteral("[[flat]]");
-        case IRInterpolationMode::NoPerspective:        return UnownedStringSlice::fromLiteral("[[center_no_perspective]]");
-        case IRInterpolationMode::Linear:               return UnownedStringSlice::fromLiteral("[[sample_no_perspective]]");
-        case IRInterpolationMode::Sample:               return UnownedStringSlice::fromLiteral("[[sample_perspective]]");
-        case IRInterpolationMode::Centroid:             return UnownedStringSlice::fromLiteral("[[center_perspective]]");
-        default:                                        return UnownedStringSlice();
+    case IRInterpolationMode::PerVertex:
+    case IRInterpolationMode::NoInterpolation: return UnownedStringSlice::fromLiteral("[[flat]]");
+    case IRInterpolationMode::NoPerspective:
+        return UnownedStringSlice::fromLiteral("[[center_no_perspective]]");
+    case IRInterpolationMode::Linear:
+        return UnownedStringSlice::fromLiteral("[[sample_no_perspective]]");
+    case IRInterpolationMode::Sample:
+        return UnownedStringSlice::fromLiteral("[[sample_perspective]]");
+    case IRInterpolationMode::Centroid:
+        return UnownedStringSlice::fromLiteral("[[center_perspective]]");
+    default: return UnownedStringSlice();
     }
 }
 
-void MetalSourceEmitter::emitInterpolationModifiersImpl(IRInst* varInst, IRType* valueType, IRVarLayout* layout)
+void MetalSourceEmitter::emitInterpolationModifiersImpl(
+    IRInst* varInst,
+    IRType* valueType,
+    IRVarLayout* layout)
 {
     SLANG_UNUSED(layout);
     SLANG_UNUSED(valueType);
@@ -1377,7 +1375,7 @@ void MetalSourceEmitter::emitInterpolationModifiersImpl(IRInst* varInst, IRType*
             continue;
 
         auto decoration = (IRInterpolationModeDecoration*)dd;
-  
+
         UnownedStringSlice modeText = _getInterpolationModifierText(decoration->getMode());
         if (modeText.getLength() > 0)
         {
@@ -1387,7 +1385,10 @@ void MetalSourceEmitter::emitInterpolationModifiersImpl(IRInst* varInst, IRType*
     }
 }
 
-void MetalSourceEmitter::emitPackOffsetModifier(IRInst* varInst, IRType* valueType, IRPackOffsetDecoration* layout)
+void MetalSourceEmitter::emitPackOffsetModifier(
+    IRInst* varInst,
+    IRType* valueType,
+    IRPackOffsetDecoration* layout)
 {
     SLANG_UNUSED(varInst);
     SLANG_UNUSED(valueType);
@@ -1395,7 +1396,9 @@ void MetalSourceEmitter::emitPackOffsetModifier(IRInst* varInst, IRType* valueTy
     // We emit packoffset as a semantic in `emitSemantic`, so nothing to do here.
 }
 
-void MetalSourceEmitter::emitRateQualifiersAndAddressSpaceImpl(IRRate* rate, AddressSpace addressSpace)
+void MetalSourceEmitter::emitRateQualifiersAndAddressSpaceImpl(
+    IRRate* rate,
+    AddressSpace addressSpace)
 {
     if (as<IRGroupSharedRate>(rate))
     {
@@ -1405,23 +1408,12 @@ void MetalSourceEmitter::emitRateQualifiersAndAddressSpaceImpl(IRRate* rate, Add
 
     switch (addressSpace)
     {
-    case AddressSpace::GroupShared:
-        m_writer->emit("threadgroup ");
-        break;
-    case AddressSpace::Uniform:
-        m_writer->emit("constant ");
-        break;
-    case AddressSpace::Global:
-        m_writer->emit("device ");
-        break;
-    case AddressSpace::ThreadLocal:
-        m_writer->emit("thread ");
-        break;
-    case AddressSpace::MetalObjectData:
-        m_writer->emit("object_data ");
-        break;
-    default:
-        break;
+    case AddressSpace::GroupShared:     m_writer->emit("threadgroup "); break;
+    case AddressSpace::Uniform:         m_writer->emit("constant "); break;
+    case AddressSpace::Global:          m_writer->emit("device "); break;
+    case AddressSpace::ThreadLocal:     m_writer->emit("thread "); break;
+    case AddressSpace::MetalObjectData: m_writer->emit("object_data "); break;
+    default:                            break;
     }
 }