format

* format

* Minor test fixes

* enable checking cpp format in ci

1 files changed, 527 insertions, 585 deletions

diff --git a/source/slang/slang-emit-wgsl.cpp b/source/slang/slang-emit-wgsl.cpp
index 105f5c3cf..ba89222cd 100644
--- a/source/slang/slang-emit-wgsl.cpp
+++ b/source/slang/slang-emit-wgsl.cpp
@@ -27,7 +27,7 @@ namespace Slang
 {
 
 void WGSLSourceEmitter::emitSwitchCaseSelectorsImpl(
-    const SwitchRegion::Case *const currentCase,
+    const SwitchRegion::Case* const currentCase,
     const bool isDefault)
 {
     // WGSL has special syntax for blocks sharing case labels:
@@ -94,9 +94,7 @@ void WGSLSourceEmitter::emitParameterGroupImpl(
             m_writer->emit(attr->getSpace());
             m_writer->emit(") ");
             break;
-
         }
-
     }
 
     auto elementType = type->getElementType();
@@ -116,43 +114,37 @@ void WGSLSourceEmitter::emitEntryPointAttributesImpl(
     switch (stage)
     {
 
-    case Stage::Fragment:
-        m_writer->emit("@fragment\n");
-        break;
-    case Stage::Vertex:
-        m_writer->emit("@vertex\n");
-        break;
+    case Stage::Fragment: m_writer->emit("@fragment\n"); break;
+    case Stage::Vertex:   m_writer->emit("@vertex\n"); break;
 
     case Stage::Compute:
-    {
-        m_writer->emit("@compute\n");
-
         {
-            Int sizeAlongAxis[kThreadGroupAxisCount];
-            getComputeThreadGroupSize(irFunc, sizeAlongAxis);
+            m_writer->emit("@compute\n");
 
-            m_writer->emit("@workgroup_size(");
-            for (int ii = 0; ii < kThreadGroupAxisCount; ++ii)
             {
-                if (ii != 0)
-                    m_writer->emit(", ");
-                m_writer->emit(sizeAlongAxis[ii]);
+                Int sizeAlongAxis[kThreadGroupAxisCount];
+                getComputeThreadGroupSize(irFunc, sizeAlongAxis);
+
+                m_writer->emit("@workgroup_size(");
+                for (int ii = 0; ii < kThreadGroupAxisCount; ++ii)
+                {
+                    if (ii != 0)
+                        m_writer->emit(", ");
+                    m_writer->emit(sizeAlongAxis[ii]);
+                }
+                m_writer->emit(")\n");
             }
-            m_writer->emit(")\n");
         }
-    }
-    break;
+        break;
 
-    default:
-        SLANG_ABORT_COMPILATION("unsupported stage.");
+    default: SLANG_ABORT_COMPILATION("unsupported stage.");
     }
-
 }
 
 // This is 'function_header' from the WGSL specification
 void WGSLSourceEmitter::emitFuncHeaderImpl(IRFunc* func)
 {
-    Slang::IRType * resultType = func->getResultType();
+    Slang::IRType* resultType = func->getResultType();
     auto name = getName(func);
 
     m_writer->emit("fn ");
@@ -182,9 +174,9 @@ void WGSLSourceEmitter::emitSimpleFuncParamImpl(IRParam* param)
 }
 
 void WGSLSourceEmitter::emitMatrixType(
-    IRType *const elementType, const IRIntegerValue& rowCountWGSL,
-    const IRIntegerValue& colCountWGSL
-    )
+    IRType* const elementType,
+    const IRIntegerValue& rowCountWGSL,
+    const IRIntegerValue& colCountWGSL)
 {
     // WGSL uses CxR convention
     m_writer->emit("mat");
@@ -206,28 +198,24 @@ static bool isPowerOf2(const uint32_t n)
     return (n != 0U) && ((n - 1U) & n) == 0U;
 }
 
-void WGSLSourceEmitter::emitStructFieldAttributes(IRStructType * structType, IRStructField * field)
+void WGSLSourceEmitter::emitStructFieldAttributes(IRStructType* structType, IRStructField* field)
 {
     // Tint emits errors unless we explicitly spell out the layout in some cases, so emit
     // offset and align attribtues for all fields.
-    IRSizeAndAlignmentDecoration *const sizeAndAlignmentDecoration =
+    IRSizeAndAlignmentDecoration* const sizeAndAlignmentDecoration =
         structType->findDecoration<IRSizeAndAlignmentDecoration>();
     // NullDifferential struct doesn't have size and alignment decoration
     if (sizeAndAlignmentDecoration == nullptr)
         return;
     SLANG_ASSERT(sizeAndAlignmentDecoration->getAlignment() > IRIntegerValue{0});
-    SLANG_ASSERT(
-        sizeAndAlignmentDecoration->getAlignment() <= IRIntegerValue{UINT32_MAX}
-    );
+    SLANG_ASSERT(sizeAndAlignmentDecoration->getAlignment() <= IRIntegerValue{UINT32_MAX});
     const uint32_t structAlignment =
         static_cast<uint32_t>(sizeAndAlignmentDecoration->getAlignment());
-    IROffsetDecoration *const fieldOffsetDecoration =
-        field->findDecoration<IROffsetDecoration>();
+    IROffsetDecoration* const fieldOffsetDecoration = field->findDecoration<IROffsetDecoration>();
     SLANG_ASSERT(fieldOffsetDecoration->getOffset() >= IRIntegerValue{0});
     SLANG_ASSERT(fieldOffsetDecoration->getOffset() <= IRIntegerValue{UINT32_MAX});
     SLANG_ASSERT(isPowerOf2(structAlignment));
-    const uint32_t fieldOffset =
-        static_cast<uint32_t>(fieldOffsetDecoration->getOffset());
+    const uint32_t fieldOffset = static_cast<uint32_t>(fieldOffsetDecoration->getOffset());
     // Alignment is GCD(fieldOffset, structAlignment)
     // TODO: Use builtin/intrinsic (e.g. __builtin_ffs)
     uint32_t fieldAlignment = 1U;
@@ -243,25 +231,15 @@ void WGSLSourceEmitter::emit(const AddressSpace addressSpace)
 {
     switch (addressSpace)
     {
-    case AddressSpace::Uniform:
-        m_writer->emit("uniform");
-        break;
+    case AddressSpace::Uniform: m_writer->emit("uniform"); break;
 
-    case AddressSpace::StorageBuffer:
-        m_writer->emit("storage");
-        break;
+    case AddressSpace::StorageBuffer: m_writer->emit("storage"); break;
 
-    case AddressSpace::Generic:
-        m_writer->emit("function");
-        break;
+    case AddressSpace::Generic: m_writer->emit("function"); break;
 
-    case AddressSpace::ThreadLocal:
-        m_writer->emit("private");
-        break;
+    case AddressSpace::ThreadLocal: m_writer->emit("private"); break;
 
-    case AddressSpace::GroupShared:
-        m_writer->emit("workgroup");
-        break;
+    case AddressSpace::GroupShared: m_writer->emit("workgroup"); break;
     }
 }
 
@@ -270,7 +248,8 @@ static const char* getWgslImageFormat(IRTextureTypeBase* type)
     // You can find the supported WGSL texel format from the URL:
     // https://www.w3.org/TR/WGSL/#storage-texel-formats
     //
-    ImageFormat imageFormat = type->hasFormat() ? (ImageFormat)type->getFormat() : ImageFormat::unknown;
+    ImageFormat imageFormat =
+        type->hasFormat() ? (ImageFormat)type->getFormat() : ImageFormat::unknown;
     switch (imageFormat)
     {
     case ImageFormat::rgba8:       return "rgba8unorm";
@@ -306,41 +285,39 @@ void WGSLSourceEmitter::emitSimpleTypeImpl(IRType* type)
     case kIROp_HLSLRWStructuredBufferType:
     case kIROp_HLSLStructuredBufferType:
     case kIROp_HLSLRasterizerOrderedStructuredBufferType:
-    {
-        auto structuredBufferType = as<IRHLSLStructuredBufferTypeBase>(type);
-        m_writer->emit("array");
-        m_writer->emit("<");
-        emitType(structuredBufferType->getElementType());
-        m_writer->emit(">");
-    }
-    break;
+        {
+            auto structuredBufferType = as<IRHLSLStructuredBufferTypeBase>(type);
+            m_writer->emit("array");
+            m_writer->emit("<");
+            emitType(structuredBufferType->getElementType());
+            m_writer->emit(">");
+        }
+        break;
 
     case kIROp_HLSLByteAddressBufferType:
     case kIROp_HLSLRWByteAddressBufferType:
-    {
-        m_writer->emit("array<u32>");
-    }
-    break;
+        {
+            m_writer->emit("array<u32>");
+        }
+        break;
 
     case kIROp_VoidType:
-    {
-        // There is no void type in WGSL.
-        // A return type of "void" is expressed by skipping the end part of the
-        // 'function_header' term:
-        // "
-        // function_header :
-        //   'fn' ident '(' param_list ? ')'
-        //       ( '->' attribute * template_elaborated_ident ) ?
-        // "
-        // In other words, in WGSL we should never even get to the point where we're
-        // asking to emit 'void'.
-        SLANG_UNEXPECTED("'void' type emitted");
-        return;
-    }
+        {
+            // There is no void type in WGSL.
+            // A return type of "void" is expressed by skipping the end part of the
+            // 'function_header' term:
+            // "
+            // function_header :
+            //   'fn' ident '(' param_list ? ')'
+            //       ( '->' attribute * template_elaborated_ident ) ?
+            // "
+            // In other words, in WGSL we should never even get to the point where we're
+            // asking to emit 'void'.
+            SLANG_UNEXPECTED("'void' type emitted");
+            return;
+        }
 
-    case kIROp_FloatType:
-        m_writer->emit("f32");
-        break;
+    case kIROp_FloatType: m_writer->emit("f32"); break;
     case kIROp_DoubleType:
         // There is no "f64" type in WGSL
         SLANG_UNEXPECTED("'double' type emitted");
@@ -354,100 +331,82 @@ void WGSLSourceEmitter::emitSimpleTypeImpl(IRType* type)
         m_f16ExtensionEnabled = true;
         m_writer->emit("f16");
         break;
-    case kIROp_BoolType:
-        m_writer->emit("bool");
-        break;
-    case kIROp_IntType:
-        m_writer->emit("i32");
-        break;
-    case kIROp_UIntType:
-        m_writer->emit("u32");
-        break;
+    case kIROp_BoolType: m_writer->emit("bool"); break;
+    case kIROp_IntType:  m_writer->emit("i32"); break;
+    case kIROp_UIntType: m_writer->emit("u32"); break;
     case kIROp_UInt64Type:
-    {
-        m_writer->emit(getDefaultBuiltinTypeName(type->getOp()));
-        return;
-    }
+        {
+            m_writer->emit(getDefaultBuiltinTypeName(type->getOp()));
+            return;
+        }
     case kIROp_Int16Type:
-    case kIROp_UInt16Type:
-        SLANG_UNEXPECTED("16 bit integer value emitted");
-        return;
+    case kIROp_UInt16Type:  SLANG_UNEXPECTED("16 bit integer value emitted"); return;
     case kIROp_Int64Type:
-    case kIROp_IntPtrType:
-        m_writer->emit("i64");
-        return;
-    case kIROp_UIntPtrType:
-        m_writer->emit("u64");
-        return;
-    case kIROp_StructType:
-        m_writer->emit(getName(type));
-        return;
+    case kIROp_IntPtrType:  m_writer->emit("i64"); return;
+    case kIROp_UIntPtrType: m_writer->emit("u64"); return;
+    case kIROp_StructType:  m_writer->emit(getName(type)); return;
 
     case kIROp_VectorType:
-    {
-        auto vecType = (IRVectorType*)type;
-        emitVectorTypeNameImpl(
-            vecType->getElementType(), getIntVal(vecType->getElementCount())
-        );
-        return;
-    }
+        {
+            auto vecType = (IRVectorType*)type;
+            emitVectorTypeNameImpl(
+                vecType->getElementType(),
+                getIntVal(vecType->getElementCount()));
+            return;
+        }
     case kIROp_MatrixType:
-    {
-        auto matType = (IRMatrixType*)type;
-        // We map matrices in Slang to WGSL matrices that represent the transpose.
-        // (See note on "terminology reversal".)
-        const IRIntegerValue colCountWGSL = getIntVal(matType->getRowCount());
-        const IRIntegerValue rowCountWGSL = getIntVal(matType->getColumnCount());
-        emitMatrixType(matType->getElementType(), rowCountWGSL, colCountWGSL);
-        return;
-    }
+        {
+            auto matType = (IRMatrixType*)type;
+            // We map matrices in Slang to WGSL matrices that represent the transpose.
+            // (See note on "terminology reversal".)
+            const IRIntegerValue colCountWGSL = getIntVal(matType->getRowCount());
+            const IRIntegerValue rowCountWGSL = getIntVal(matType->getColumnCount());
+            emitMatrixType(matType->getElementType(), rowCountWGSL, colCountWGSL);
+            return;
+        }
     case kIROp_SamplerStateType:
-    {
-        m_writer->emit("sampler");
-        return;
-    }
+        {
+            m_writer->emit("sampler");
+            return;
+        }
 
     case kIROp_SamplerComparisonStateType:
-    {
-        m_writer->emit("sampler_comparison");
-        return;
-    }
+        {
+            m_writer->emit("sampler_comparison");
+            return;
+        }
 
     case kIROp_PtrType:
     case kIROp_InOutType:
     case kIROp_OutType:
     case kIROp_RefType:
     case kIROp_ConstRefType:
-    {
-        auto ptrType = cast<IRPtrTypeBase>(type);
-        m_writer->emit("ptr<");
-        emit((AddressSpace)ptrType->getAddressSpace());
-        m_writer->emit(", ");
-        emitType((IRType*)ptrType->getValueType());
-        m_writer->emit(">");
-        return;
-    }
+        {
+            auto ptrType = cast<IRPtrTypeBase>(type);
+            m_writer->emit("ptr<");
+            emit((AddressSpace)ptrType->getAddressSpace());
+            m_writer->emit(", ");
+            emitType((IRType*)ptrType->getValueType());
+            m_writer->emit(">");
+            return;
+        }
 
     case kIROp_ArrayType:
-    {
-        m_writer->emit("array<");
-        emitType((IRType*)type->getOperand(0));
-        m_writer->emit(", ");
-        emitVal(type->getOperand(1), getInfo(EmitOp::General));
-        m_writer->emit(">");
-        return;
-    }
+        {
+            m_writer->emit("array<");
+            emitType((IRType*)type->getOperand(0));
+            m_writer->emit(", ");
+            emitVal(type->getOperand(1), getInfo(EmitOp::General));
+            m_writer->emit(">");
+            return;
+        }
     case kIROp_TextureType:
         if (auto texType = as<IRTextureType>(type))
         {
             switch (texType->getAccess())
             {
-            case SLANG_RESOURCE_ACCESS_READ_WRITE:
-                m_writer->emit("texture_storage");
-                break;
-            default:
-                m_writer->emit("texture");
-                break;
+            case SLANG_RESOURCE_ACCESS_READ_WRITE: m_writer->emit("texture_storage"); break;
+            default:                               m_writer->emit("texture"); break;
             }
 
             if (texType->isShadow())
@@ -462,9 +421,9 @@ void WGSLSourceEmitter::emitSimpleTypeImpl(IRType* type)
 
             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_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;
             }
 
@@ -498,17 +457,14 @@ void WGSLSourceEmitter::emitSimpleTypeImpl(IRType* type)
         return;
 
     case kIROp_AtomicType:
-    {
-        m_writer->emit("atomic<");
-        emitType(cast<IRAtomicType>(type)->getElementType());
-        m_writer->emit(">");
-        return;
-    }
-    default:
-        break;
-
+        {
+            m_writer->emit("atomic<");
+            emitType(cast<IRAtomicType>(type)->getElementType());
+            m_writer->emit(">");
+            return;
+        }
+    default: break;
     }
-
 }
 
 void WGSLSourceEmitter::emitLayoutQualifiersImpl(IRVarLayout* layout)
@@ -535,18 +491,15 @@ void WGSLSourceEmitter::emitLayoutQualifiersImpl(IRVarLayout* layout)
             return;
         }
     }
-
 }
 
-void WGSLSourceEmitter::emitVarKeywordImpl(IRType * type, IRInst* varDecl)
+void WGSLSourceEmitter::emitVarKeywordImpl(IRType* type, IRInst* varDecl)
 {
     switch (varDecl->getOp())
     {
     case kIROp_GlobalParam:
     case kIROp_GlobalVar:
-    case kIROp_Var:
-        m_writer->emit("var");
-        break;
+    case kIROp_Var:         m_writer->emit("var"); break;
     default:
         if (as<IRModuleInst>(varDecl->getParent()))
         {
@@ -563,7 +516,8 @@ void WGSLSourceEmitter::emitVarKeywordImpl(IRType * type, IRInst* varDecl)
     {
         m_writer->emit("<workgroup>");
     }
-    else if (type->getOp() == kIROp_HLSLRWStructuredBufferType ||
+    else if (
+        type->getOp() == kIROp_HLSLRWStructuredBufferType ||
         type->getOp() == kIROp_HLSLRasterizerOrderedStructuredBufferType ||
         type->getOp() == kIROp_HLSLRWByteAddressBufferType)
     {
@@ -571,14 +525,15 @@ void WGSLSourceEmitter::emitVarKeywordImpl(IRType * type, IRInst* varDecl)
         m_writer->emit("storage, read_write");
         m_writer->emit(">");
     }
-    else if (type->getOp() == kIROp_HLSLStructuredBufferType ||
+    else if (
+        type->getOp() == kIROp_HLSLStructuredBufferType ||
         type->getOp() == kIROp_HLSLByteAddressBufferType)
     {
         m_writer->emit("<");
         m_writer->emit("storage, read");
         m_writer->emit(">");
     }
-    else if(varDecl->getOp() == kIROp_GlobalVar)
+    else if (varDecl->getOp() == kIROp_GlobalVar)
     {
         // Global ("module-scope") non-handle variables need to specify storage space
 
@@ -604,90 +559,85 @@ void WGSLSourceEmitter::_emitType(IRType* type, DeclaratorInfo* declarator)
     {
     case kIROp_ArrayType:
     case kIROp_AttributedType:
-    case kIROp_UnsizedArrayType:
-        emitSimpleTypeAndDeclarator(type, declarator);
-        break;
-    default:
-        CLikeSourceEmitter::_emitType(type, declarator);
-        break;
+    case kIROp_UnsizedArrayType: emitSimpleTypeAndDeclarator(type, declarator); break;
+    default:                     CLikeSourceEmitter::_emitType(type, declarator); break;
     }
 }
 
 void WGSLSourceEmitter::emitDeclaratorImpl(DeclaratorInfo* declarator)
 {
-    if (!declarator) return;
+    if (!declarator)
+        return;
 
     m_writer->emit(" ");
 
     switch (declarator->flavor)
     {
     case DeclaratorInfo::Flavor::Name:
-    {
-        auto nameDeclarator = (NameDeclaratorInfo*)declarator;
-        m_writer->emitName(*nameDeclarator->nameAndLoc);
-    }
-    break;
+        {
+            auto nameDeclarator = (NameDeclaratorInfo*)declarator;
+            m_writer->emitName(*nameDeclarator->nameAndLoc);
+        }
+        break;
 
     case DeclaratorInfo::Flavor::SizedArray:
-    {
-        // Sized arrays are just types (array<T, N>) in WGSL -- they are not
-        // supported at the syntax level
-        // https://www.w3.org/TR/WGSL/#array
-        SLANG_UNEXPECTED("Sized array declarator");
-    }
-    break;
+        {
+            // Sized arrays are just types (array<T, N>) in WGSL -- they are not
+            // supported at the syntax level
+            // https://www.w3.org/TR/WGSL/#array
+            SLANG_UNEXPECTED("Sized array declarator");
+        }
+        break;
 
     case DeclaratorInfo::Flavor::UnsizedArray:
-    {
-        // Unsized arrays are just types (array<T>) in WGSL -- they are not
-        // supported at the syntax level
-        // https://www.w3.org/TR/WGSL/#array
-        SLANG_UNEXPECTED("Unsized array declarator");
-    }
-    break;
+        {
+            // Unsized arrays are just types (array<T>) in WGSL -- they are not
+            // supported at the syntax level
+            // https://www.w3.org/TR/WGSL/#array
+            SLANG_UNEXPECTED("Unsized array declarator");
+        }
+        break;
 
     case DeclaratorInfo::Flavor::Ptr:
-    {
-        // Pointers (ptr<AS,T,AM>) are just types in WGSL -- they are not supported at
-        // the syntax level
-        // https://www.w3.org/TR/WGSL/#ref-ptr-types
-        SLANG_UNEXPECTED("Pointer declarator");
-    }
-    break;
+        {
+            // Pointers (ptr<AS,T,AM>) are just types in WGSL -- they are not supported at
+            // the syntax level
+            // https://www.w3.org/TR/WGSL/#ref-ptr-types
+            SLANG_UNEXPECTED("Pointer declarator");
+        }
+        break;
 
     case DeclaratorInfo::Flavor::Ref:
-    {
-        // References (ref<AS,T,AM>) are just types in WGSL -- they are not supported
-        // at the syntax level
-        // https://www.w3.org/TR/WGSL/#ref-ptr-types
-        SLANG_UNEXPECTED("Reference declarator");
-    }
-    break;
+        {
+            // References (ref<AS,T,AM>) are just types in WGSL -- they are not supported
+            // at the syntax level
+            // https://www.w3.org/TR/WGSL/#ref-ptr-types
+            SLANG_UNEXPECTED("Reference declarator");
+        }
+        break;
 
     case DeclaratorInfo::Flavor::LiteralSizedArray:
-    {
-        // Sized arrays are just types (array<T, N>) in WGSL -- they are not supported
-        // at the syntax level
-        // https://www.w3.org/TR/WGSL/#array
-        SLANG_UNEXPECTED("Literal-sized array declarator");
-    }
-    break;
+        {
+            // Sized arrays are just types (array<T, N>) in WGSL -- they are not supported
+            // at the syntax level
+            // https://www.w3.org/TR/WGSL/#array
+            SLANG_UNEXPECTED("Literal-sized array declarator");
+        }
+        break;
 
     case DeclaratorInfo::Flavor::Attributed:
-    {
-        auto attributedDeclarator = (AttributedDeclaratorInfo*)declarator;
-        auto instWithAttributes = attributedDeclarator->instWithAttributes;
-        for (auto attr : instWithAttributes->getAllAttrs())
         {
-            _emitPostfixTypeAttr(attr);
+            auto attributedDeclarator = (AttributedDeclaratorInfo*)declarator;
+            auto instWithAttributes = attributedDeclarator->instWithAttributes;
+            for (auto attr : instWithAttributes->getAllAttrs())
+            {
+                _emitPostfixTypeAttr(attr);
+            }
+            emitDeclarator(attributedDeclarator->next);
         }
-        emitDeclarator(attributedDeclarator->next);
-    }
-    break;
-
-    default:
-        SLANG_DIAGNOSE_UNEXPECTED(getSink(), SourceLoc(), "unknown declarator flavor");
         break;
+
+    default: SLANG_DIAGNOSE_UNEXPECTED(getSink(), SourceLoc(), "unknown declarator flavor"); break;
     }
 }
 
@@ -708,9 +658,7 @@ void WGSLSourceEmitter::emitOperandImpl(IRInst* operand, EmitOpInfo const& outer
     }
 }
 
-void WGSLSourceEmitter::emitSimpleTypeAndDeclaratorImpl(
-    IRType* type,
-    DeclaratorInfo* declarator)
+void WGSLSourceEmitter::emitSimpleTypeAndDeclaratorImpl(IRType* type, DeclaratorInfo* declarator)
 {
     if (declarator)
     {
@@ -725,151 +673,145 @@ void WGSLSourceEmitter::emitSimpleValueImpl(IRInst* inst)
     switch (inst->getOp())
     {
     case kIROp_IntLit:
-    {
-        auto litInst = static_cast<IRConstant*>(inst);
-
-        IRBasicType* type = as<IRBasicType>(inst->getDataType());
-        if (type)
         {
-            switch (type->getBaseType())
+            auto litInst = static_cast<IRConstant*>(inst);
+
+            IRBasicType* type = as<IRBasicType>(inst->getDataType());
+            if (type)
             {
+                switch (type->getBaseType())
+                {
                 default:
 
                 case BaseType::Int8:
                 case BaseType::UInt8:
-                {
-                    SLANG_UNEXPECTED("8 bit integer value emitted");
-                    break;
-                }
+                    {
+                        SLANG_UNEXPECTED("8 bit integer value emitted");
+                        break;
+                    }
                 case BaseType::Int16:
                 case BaseType::UInt16:
-                {
-                    SLANG_UNEXPECTED("16 bit integer value emitted");
-                    break;
-                }
+                    {
+                        SLANG_UNEXPECTED("16 bit integer value emitted");
+                        break;
+                    }
                 case BaseType::Int:
-                {
-                    m_writer->emit("i32(");
-                    m_writer->emit(int32_t(litInst->value.intVal));
-                    m_writer->emit(")");
-                    return;
-                }
+                    {
+                        m_writer->emit("i32(");
+                        m_writer->emit(int32_t(litInst->value.intVal));
+                        m_writer->emit(")");
+                        return;
+                    }
                 case BaseType::UInt:
-                {
-                    m_writer->emit("u32(");
-                    m_writer->emit(UInt(uint32_t(litInst->value.intVal)));
-                    m_writer->emit(")");
-                    break;
-                }
+                    {
+                        m_writer->emit("u32(");
+                        m_writer->emit(UInt(uint32_t(litInst->value.intVal)));
+                        m_writer->emit(")");
+                        break;
+                    }
                 case BaseType::Int64:
-                {
-                    m_writer->emit("i64(");
-                    m_writer->emitInt64(int64_t(litInst->value.intVal));
-                    m_writer->emit(")");
-                    break;
-                }
+                    {
+                        m_writer->emit("i64(");
+                        m_writer->emitInt64(int64_t(litInst->value.intVal));
+                        m_writer->emit(")");
+                        break;
+                    }
                 case BaseType::UInt64:
-                {
-                    m_writer->emit("u64(");
-                    SLANG_COMPILE_TIME_ASSERT(
-                        sizeof(litInst->value.intVal) >= sizeof(uint64_t)
-                    );
-                    m_writer->emitUInt64(uint64_t(litInst->value.intVal));
-                    m_writer->emit(")");
-                    break;
-                }
+                    {
+                        m_writer->emit("u64(");
+                        SLANG_COMPILE_TIME_ASSERT(
+                            sizeof(litInst->value.intVal) >= sizeof(uint64_t));
+                        m_writer->emitUInt64(uint64_t(litInst->value.intVal));
+                        m_writer->emit(")");
+                        break;
+                    }
                 case BaseType::IntPtr:
-                {
+                    {
 #if SLANG_PTR_IS_64
-                    m_writer->emit("i64(");
-                    m_writer->emitInt64(int64_t(litInst->value.intVal));
-                    m_writer->emit(")");
+                        m_writer->emit("i64(");
+                        m_writer->emitInt64(int64_t(litInst->value.intVal));
+                        m_writer->emit(")");
 #else
-                    m_writer->emit("i32(");
-                    m_writer->emit(int(litInst->value.intVal));
-                    m_writer->emit(")");
+                        m_writer->emit("i32(");
+                        m_writer->emit(int(litInst->value.intVal));
+                        m_writer->emit(")");
 #endif
-                    break;
-                }
+                        break;
+                    }
                 case BaseType::UIntPtr:
-                {
+                    {
 #if SLANG_PTR_IS_64
-                    m_writer->emit("u64(");
-                    m_writer->emitUInt64(uint64_t(litInst->value.intVal));
-                    m_writer->emit(")");
+                        m_writer->emit("u64(");
+                        m_writer->emitUInt64(uint64_t(litInst->value.intVal));
+                        m_writer->emit(")");
 #else
-                    m_writer->emit("u32(");
-                    m_writer->emit(UInt(uint32_t(litInst->value.intVal)));
-                    m_writer->emit(")");
+                        m_writer->emit("u32(");
+                        m_writer->emit(UInt(uint32_t(litInst->value.intVal)));
+                        m_writer->emit(")");
 #endif
-                    break;
+                        break;
+                    }
                 }
-
             }
+            else
+            {
+                // If no type... just output what we have
+                m_writer->emit(litInst->value.intVal);
+            }
+            break;
         }
-        else
-        {
-            // If no type... just output what we have
-            m_writer->emit(litInst->value.intVal);
-        }
-        break;
-    }
 
     case kIROp_FloatLit:
-    {
-        auto litInst = static_cast<IRConstant*>(inst);
-
-        IRBasicType* type = as<IRBasicType>(inst->getDataType());
-        if (type)
         {
-            switch (type->getBaseType())
-            {
-            default:
+            auto litInst = static_cast<IRConstant*>(inst);
 
-            case BaseType::Half:
+            IRBasicType* type = as<IRBasicType>(inst->getDataType());
+            if (type)
             {
-                m_writer->emit(litInst->value.floatVal);
-                m_writer->emit("h");
-                m_f16ExtensionEnabled = true;
-            }
-            break;
+                switch (type->getBaseType())
+                {
+                default:
 
-            case BaseType::Float:
-            {
-                m_writer->emit(litInst->value.floatVal);
-                m_writer->emit("f");
-            }
-            break;
+                case BaseType::Half:
+                    {
+                        m_writer->emit(litInst->value.floatVal);
+                        m_writer->emit("h");
+                        m_f16ExtensionEnabled = true;
+                    }
+                    break;
 
-            case BaseType::Double:
-            {
-                 // There is not "f64" in WGSL
-                SLANG_UNEXPECTED("'double' type emitted");
+                case BaseType::Float:
+                    {
+                        m_writer->emit(litInst->value.floatVal);
+                        m_writer->emit("f");
+                    }
+                    break;
+
+                case BaseType::Double:
+                    {
+                        // There is not "f64" in WGSL
+                        SLANG_UNEXPECTED("'double' type emitted");
+                    }
+                    break;
+                }
             }
-            break;
+            else
+            {
+                // If no type... just output what we have
+                m_writer->emit(litInst->value.floatVal);
             }
         }
-        else
-        {
-            // If no type... just output what we have
-            m_writer->emit(litInst->value.floatVal);
-        }
-    }
-    break;
+        break;
 
     case kIROp_BoolLit:
-    {
-        bool val = ((IRConstant*)inst)->value.intVal != 0;
-        m_writer->emit(val ? "true" : "false");
-    }
-    break;
-
-    default:
-        SLANG_UNIMPLEMENTED_X("val case for emit");
+        {
+            bool val = ((IRConstant*)inst)->value.intVal != 0;
+            m_writer->emit(val ? "true" : "false");
+        }
         break;
-    }
-
 
+    default: SLANG_UNIMPLEMENTED_X("val case for emit"); break;
+    }
 }
 
 void WGSLSourceEmitter::emitParamTypeImpl(IRType* type, const String& name)
@@ -881,137 +823,136 @@ bool WGSLSourceEmitter::tryEmitInstStmtImpl(IRInst* inst)
 {
     switch (inst->getOp())
     {
-    default:
-        return false;
+    default: return false;
     case kIROp_AtomicLoad:
-    {
-        emitInstResultDecl(inst);
-        m_writer->emit("atomicLoad(&(");
-        emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
-        m_writer->emit("));\n");
-        return true;
-    }
+        {
+            emitInstResultDecl(inst);
+            m_writer->emit("atomicLoad(&(");
+            emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
+            m_writer->emit("));\n");
+            return true;
+        }
     case kIROp_AtomicStore:
-    {
-        m_writer->emit("atomicStore(&(");
-        emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
-        m_writer->emit("), ");
-        emitOperand(inst->getOperand(1), getInfo(EmitOp::General));
-        m_writer->emit(");\n");
-        return true;
-    }
+        {
+            m_writer->emit("atomicStore(&(");
+            emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
+            m_writer->emit("), ");
+            emitOperand(inst->getOperand(1), getInfo(EmitOp::General));
+            m_writer->emit(");\n");
+            return true;
+        }
     case kIROp_AtomicExchange:
-    {
-        emitInstResultDecl(inst);
-        m_writer->emit("atomicExchange(&(");
-        emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
-        m_writer->emit("), ");
-        emitOperand(inst->getOperand(1), getInfo(EmitOp::General));
-        m_writer->emit(");\n");
-        return true;
-    }
+        {
+            emitInstResultDecl(inst);
+            m_writer->emit("atomicExchange(&(");
+            emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
+            m_writer->emit("), ");
+            emitOperand(inst->getOperand(1), getInfo(EmitOp::General));
+            m_writer->emit(");\n");
+            return true;
+        }
     case kIROp_AtomicCompareExchange:
-    {
-        emitInstResultDecl(inst);
-        m_writer->emit("atomicCompareExchangeWeak(&(");
-        emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
-        m_writer->emit("), ");
-        emitOperand(inst->getOperand(1), getInfo(EmitOp::General));
-        m_writer->emit(", ");
-        emitOperand(inst->getOperand(2), getInfo(EmitOp::General));
-        m_writer->emit(").old_value;\n");
-        return true;
-    }
+        {
+            emitInstResultDecl(inst);
+            m_writer->emit("atomicCompareExchangeWeak(&(");
+            emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
+            m_writer->emit("), ");
+            emitOperand(inst->getOperand(1), getInfo(EmitOp::General));
+            m_writer->emit(", ");
+            emitOperand(inst->getOperand(2), getInfo(EmitOp::General));
+            m_writer->emit(").old_value;\n");
+            return true;
+        }
     case kIROp_AtomicAdd:
-    {
-        emitInstResultDecl(inst);
-        m_writer->emit("atomicAdd(&(");
-        emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
-        m_writer->emit("), ");
-        emitOperand(inst->getOperand(1), getInfo(EmitOp::General));
-        m_writer->emit(");\n");
-        return true;
-    }
+        {
+            emitInstResultDecl(inst);
+            m_writer->emit("atomicAdd(&(");
+            emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
+            m_writer->emit("), ");
+            emitOperand(inst->getOperand(1), getInfo(EmitOp::General));
+            m_writer->emit(");\n");
+            return true;
+        }
     case kIROp_AtomicSub:
-    {
-        emitInstResultDecl(inst);
-        m_writer->emit("atomicSub(&(");
-        emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
-        m_writer->emit("), ");
-        emitOperand(inst->getOperand(1), getInfo(EmitOp::General));
-        m_writer->emit(");\n");
-        return true;
-    }
+        {
+            emitInstResultDecl(inst);
+            m_writer->emit("atomicSub(&(");
+            emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
+            m_writer->emit("), ");
+            emitOperand(inst->getOperand(1), getInfo(EmitOp::General));
+            m_writer->emit(");\n");
+            return true;
+        }
     case kIROp_AtomicAnd:
-    {
-        emitInstResultDecl(inst);
-        m_writer->emit("atomicAnd(&(");
-        emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
-        m_writer->emit("), ");
-        emitOperand(inst->getOperand(1), getInfo(EmitOp::General));
-        m_writer->emit(");\n");
-        return true;
-    }
+        {
+            emitInstResultDecl(inst);
+            m_writer->emit("atomicAnd(&(");
+            emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
+            m_writer->emit("), ");
+            emitOperand(inst->getOperand(1), getInfo(EmitOp::General));
+            m_writer->emit(");\n");
+            return true;
+        }
     case kIROp_AtomicOr:
-    {
-        emitInstResultDecl(inst);
-        m_writer->emit("atomicOr(&(");
-        emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
-        m_writer->emit("), ");
-        emitOperand(inst->getOperand(1), getInfo(EmitOp::General));
-        m_writer->emit(");\n");
-        return true;
-    }
+        {
+            emitInstResultDecl(inst);
+            m_writer->emit("atomicOr(&(");
+            emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
+            m_writer->emit("), ");
+            emitOperand(inst->getOperand(1), getInfo(EmitOp::General));
+            m_writer->emit(");\n");
+            return true;
+        }
     case kIROp_AtomicXor:
-    {
-        emitInstResultDecl(inst);
-        m_writer->emit("atomicXor(&(");
-        emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
-        m_writer->emit("), ");
-        emitOperand(inst->getOperand(1), getInfo(EmitOp::General));
-        m_writer->emit(");\n");
-        return true;
-    }
+        {
+            emitInstResultDecl(inst);
+            m_writer->emit("atomicXor(&(");
+            emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
+            m_writer->emit("), ");
+            emitOperand(inst->getOperand(1), getInfo(EmitOp::General));
+            m_writer->emit(");\n");
+            return true;
+        }
     case kIROp_AtomicMin:
-    {
-        emitInstResultDecl(inst);
-        m_writer->emit("atomicMin(&(");
-        emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
-        m_writer->emit("), ");
-        emitOperand(inst->getOperand(1), getInfo(EmitOp::General));
-        m_writer->emit(");\n");
-        return true;
-    }
+        {
+            emitInstResultDecl(inst);
+            m_writer->emit("atomicMin(&(");
+            emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
+            m_writer->emit("), ");
+            emitOperand(inst->getOperand(1), getInfo(EmitOp::General));
+            m_writer->emit(");\n");
+            return true;
+        }
     case kIROp_AtomicMax:
-    {
-        emitInstResultDecl(inst);
-        m_writer->emit("atomicMax(&(");
-        emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
-        m_writer->emit("), ");
-        emitOperand(inst->getOperand(1), getInfo(EmitOp::General));
-        m_writer->emit(");\n");
-        return true;
-    }
+        {
+            emitInstResultDecl(inst);
+            m_writer->emit("atomicMax(&(");
+            emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
+            m_writer->emit("), ");
+            emitOperand(inst->getOperand(1), getInfo(EmitOp::General));
+            m_writer->emit(");\n");
+            return true;
+        }
     case kIROp_AtomicInc:
-    {
-        emitInstResultDecl(inst);
-        m_writer->emit("atomicAdd(&(");
-        emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
-        m_writer->emit("), ");
-        emitType(inst->getDataType());
-        m_writer->emit("(1));\n");
-        return true;
-    }
+        {
+            emitInstResultDecl(inst);
+            m_writer->emit("atomicAdd(&(");
+            emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
+            m_writer->emit("), ");
+            emitType(inst->getDataType());
+            m_writer->emit("(1));\n");
+            return true;
+        }
     case kIROp_AtomicDec:
-    {
-        emitInstResultDecl(inst);
-        m_writer->emit("atomicSub(&(");
-        emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
-        m_writer->emit("), ");
-        emitType(inst->getDataType());
-        m_writer->emit("(1));\n");
-        return true;
-    }
+        {
+            emitInstResultDecl(inst);
+            m_writer->emit("atomicSub(&(");
+            emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
+            m_writer->emit("), ");
+            emitType(inst->getDataType());
+            m_writer->emit("(1));\n");
+            return true;
+        }
     }
 }
 
@@ -1039,148 +980,148 @@ bool WGSLSourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inOu
     switch (inst->getOp())
     {
     case kIROp_MakeVectorFromScalar:
-    {
-        // In WGSL this is done by calling the vec* overloads listed in [1]
-        // [1] https://www.w3.org/TR/WGSL/#value-constructor-builtin-function
-        emitType(inst->getDataType());
-        m_writer->emit("(");
-        auto prec = getInfo(EmitOp::Prefix);
-        emitOperand(inst->getOperand(0), rightSide(outerPrec, prec));
-        m_writer->emit(")");
-        return true;
-    }
-    break;
+        {
+            // In WGSL this is done by calling the vec* overloads listed in [1]
+            // [1] https://www.w3.org/TR/WGSL/#value-constructor-builtin-function
+            emitType(inst->getDataType());
+            m_writer->emit("(");
+            auto prec = getInfo(EmitOp::Prefix);
+            emitOperand(inst->getOperand(0), rightSide(outerPrec, prec));
+            m_writer->emit(")");
+            return true;
+        }
+        break;
 
     case kIROp_BitCast:
-    {
-        // In WGSL there is a built-in bitcast function!
-        // https://www.w3.org/TR/WGSL/#bitcast-builtin
-        m_writer->emit("bitcast");
-        m_writer->emit("<");
-        emitType(inst->getDataType());
-        m_writer->emit(">");
-        m_writer->emit("(");
-        emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
-        m_writer->emit(")");
-        return true;
-    }
-    break;
+        {
+            // In WGSL there is a built-in bitcast function!
+            // https://www.w3.org/TR/WGSL/#bitcast-builtin
+            m_writer->emit("bitcast");
+            m_writer->emit("<");
+            emitType(inst->getDataType());
+            m_writer->emit(">");
+            m_writer->emit("(");
+            emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
+            m_writer->emit(")");
+            return true;
+        }
+        break;
 
     case kIROp_MakeArray:
     case kIROp_MakeStruct:
-    {
-        // It seems there are currently no designated initializers in WGSL.
-        // Similarly for array initializers.
-        // https://github.com/gpuweb/gpuweb/issues/4210
-
-        // There is a constructor named like the struct/array type itself
-        auto type = inst->getDataType();
-        emitType(type);
-        m_writer->emit("( ");
-        UInt argCount = inst->getOperandCount();
-        for (UInt aa = 0; aa < argCount; ++aa)
         {
-            if (aa != 0) m_writer->emit(", ");
-            emitOperand(inst->getOperand(aa), getInfo(EmitOp::General));
-        }
-        m_writer->emit(" )");
+            // It seems there are currently no designated initializers in WGSL.
+            // Similarly for array initializers.
+            // https://github.com/gpuweb/gpuweb/issues/4210
+
+            // There is a constructor named like the struct/array type itself
+            auto type = inst->getDataType();
+            emitType(type);
+            m_writer->emit("( ");
+            UInt argCount = inst->getOperandCount();
+            for (UInt aa = 0; aa < argCount; ++aa)
+            {
+                if (aa != 0)
+                    m_writer->emit(", ");
+                emitOperand(inst->getOperand(aa), getInfo(EmitOp::General));
+            }
+            m_writer->emit(" )");
 
-        return true;
-    }
-    break;
+            return true;
+        }
+        break;
 
     case kIROp_MakeArrayFromElement:
-    {
-        // It seems there are currently no array initializers in WGSL.
-
-        // There is a constructor named like the array type itself
-        auto type = inst->getDataType();
-        emitType(type);
-        m_writer->emit("(");
-        UInt argCount =
-            (UInt)cast<IRIntLit>(
-                cast<IRArrayType>(inst->getDataType())->getElementCount()
-            )->getValue();
-        for (UInt aa = 0; aa < argCount; ++aa)
         {
-            if (aa != 0) m_writer->emit(", ");
-            emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
+            // It seems there are currently no array initializers in WGSL.
+
+            // There is a constructor named like the array type itself
+            auto type = inst->getDataType();
+            emitType(type);
+            m_writer->emit("(");
+            UInt argCount =
+                (UInt)cast<IRIntLit>(cast<IRArrayType>(inst->getDataType())->getElementCount())
+                    ->getValue();
+            for (UInt aa = 0; aa < argCount; ++aa)
+            {
+                if (aa != 0)
+                    m_writer->emit(", ");
+                emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
+            }
+            m_writer->emit(")");
+            return true;
         }
-        m_writer->emit(")");
-        return true;
-    }
-    break;
+        break;
 
     case kIROp_StructuredBufferLoad:
     case kIROp_RWStructuredBufferLoad:
     case kIROp_RWStructuredBufferGetElementPtr:
-    {
-        emitOperand(inst->getOperand(0), leftSide(outerPrec, getInfo(EmitOp::Postfix)));
-        m_writer->emit("[");
-        emitOperand(inst->getOperand(1), getInfo(EmitOp::General));
-        m_writer->emit("]");
-        return true;
-    }
-    break;
+        {
+            emitOperand(inst->getOperand(0), leftSide(outerPrec, getInfo(EmitOp::Postfix)));
+            m_writer->emit("[");
+            emitOperand(inst->getOperand(1), getInfo(EmitOp::General));
+            m_writer->emit("]");
+            return true;
+        }
+        break;
 
     case kIROp_Rsh:
     case kIROp_Lsh:
-    {
-        // Shift amounts must be an unsigned type in WGSL
-        // https://www.w3.org/TR/WGSL/#bit-expr
-        IRInst *const shiftAmount = inst->getOperand(1);
-        IRType *const shiftAmountType = shiftAmount->getDataType();
-        if (shiftAmountType->getOp() == kIROp_IntType)
         {
-            // Dawn complains about "mixing '<<' and '|' requires parenthesis", so let's
-            // add parenthesis.
-            m_writer->emit("(");
-
-            const auto emitOp = getEmitOpForOp(inst->getOp());
-            const auto info = getInfo(emitOp);
-
-            const bool needClose = maybeEmitParens(outerPrec, info);
-            emitOperand(inst->getOperand(0), leftSide(outerPrec, info));
-            m_writer->emit(" ");
-            m_writer->emit(info.op);
-            m_writer->emit(" ");
-            m_writer->emit("bitcast<u32>(");
-            emitOperand(inst->getOperand(1), rightSide(outerPrec, info));
-            m_writer->emit(")");
-            maybeCloseParens(needClose);
-
-            m_writer->emit(")");
-            return true;
+            // Shift amounts must be an unsigned type in WGSL
+            // https://www.w3.org/TR/WGSL/#bit-expr
+            IRInst* const shiftAmount = inst->getOperand(1);
+            IRType* const shiftAmountType = shiftAmount->getDataType();
+            if (shiftAmountType->getOp() == kIROp_IntType)
+            {
+                // Dawn complains about "mixing '<<' and '|' requires parenthesis", so let's
+                // add parenthesis.
+                m_writer->emit("(");
+
+                const auto emitOp = getEmitOpForOp(inst->getOp());
+                const auto info = getInfo(emitOp);
+
+                const bool needClose = maybeEmitParens(outerPrec, info);
+                emitOperand(inst->getOperand(0), leftSide(outerPrec, info));
+                m_writer->emit(" ");
+                m_writer->emit(info.op);
+                m_writer->emit(" ");
+                m_writer->emit("bitcast<u32>(");
+                emitOperand(inst->getOperand(1), rightSide(outerPrec, info));
+                m_writer->emit(")");
+                maybeCloseParens(needClose);
+
+                m_writer->emit(")");
+                return true;
+            }
         }
-    }
-    break;
+        break;
 
     case kIROp_ByteAddressBufferLoad:
-    {
-        // Indices in Slang code count bytes, but in WASM they count u32's since
-        // byte address buffers translate to array<u32> in WASM, so divide by 4.
-        emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
-        m_writer->emit("[(");
-        emitOperand(inst->getOperand(1), getInfo(EmitOp::General));
-        m_writer->emit(")/4]");
-        return true;
-    }
-    break;
+        {
+            // Indices in Slang code count bytes, but in WASM they count u32's since
+            // byte address buffers translate to array<u32> in WASM, so divide by 4.
+            emitOperand(inst->getOperand(0), getInfo(EmitOp::General));
+            m_writer->emit("[(");
+            emitOperand(inst->getOperand(1), getInfo(EmitOp::General));
+            m_writer->emit(")/4]");
+            return true;
+        }
+        break;
 
     case kIROp_ByteAddressBufferStore:
-    {
-        // Indices in Slang code count bytes, but in WASM they count u32's since
-        // byte address buffers translate to array<u32> in WASM, so divide by 4.
-        auto base = inst->getOperand(0);
-        emitOperand(base, EmitOpInfo());
-        m_writer->emit("[(");
-        emitOperand(inst->getOperand(1), getInfo(EmitOp::General));
-        m_writer->emit(")/4] = ");
-        emitOperand(inst->getOperand(inst->getOperandCount() - 1), getInfo(EmitOp::General));
-        return true;
-    }
-    break;
-
+        {
+            // Indices in Slang code count bytes, but in WASM they count u32's since
+            // byte address buffers translate to array<u32> in WASM, so divide by 4.
+            auto base = inst->getOperand(0);
+            emitOperand(base, EmitOpInfo());
+            m_writer->emit("[(");
+            emitOperand(inst->getOperand(1), getInfo(EmitOp::General));
+            m_writer->emit(")/4] = ");
+            emitOperand(inst->getOperand(inst->getOperandCount() - 1), getInfo(EmitOp::General));
+            return true;
+        }
+        break;
     }
 
     return false;
@@ -1216,8 +1157,7 @@ void WGSLSourceEmitter::emitIntrinsicCallExprImpl(
     IRCall* inst,
     UnownedStringSlice intrinsicDefinition,
     IRInst* intrinsicInst,
-    EmitOpInfo const& inOuterPrec
-    )
+    EmitOpInfo const& inOuterPrec)
 {
     // The f16 constructor is generated for f32tof16
     if (intrinsicDefinition.startsWith("f16"))
@@ -1226,8 +1166,10 @@ void WGSLSourceEmitter::emitIntrinsicCallExprImpl(
     }
 
     CLikeSourceEmitter::emitIntrinsicCallExprImpl(
-        inst, intrinsicDefinition, intrinsicInst, inOuterPrec
-    );
+        inst,
+        intrinsicDefinition,
+        intrinsicInst,
+        inOuterPrec);
 }
 
 } // namespace Slang