[WHLSL] Make generated Metal code should be indented properly to ease reading while...
authorweinig@apple.com <weinig@apple.com@268f45cc-cd09-0410-ab3c-d52691b4dbfc>
Tue, 20 Aug 2019 05:47:38 +0000 (05:47 +0000)
committerweinig@apple.com <weinig@apple.com@268f45cc-cd09-0410-ab3c-d52691b4dbfc>
Tue, 20 Aug 2019 05:47:38 +0000 (05:47 +0000)
https://bugs.webkit.org/show_bug.cgi?id=200870

Reviewed by Saam Barati.

Source/WebCore:

Adds indentation of generated Metal code using the new WTF::Indentation and WTF::IndentationScope
classes. This makes reading the generated Metal code much easier while debugging it, and does not
seem to introduce any performance regression (I measured no difference on compute_boids.html test).

* Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp:
(WebCore::WHLSL::Metal::EntryPointScaffolding::emitResourceHelperTypes):
(WebCore::WHLSL::Metal::EntryPointScaffolding::emitUnpackResourcesAndNamedBuiltIns):
(WebCore::WHLSL::Metal::VertexEntryPointScaffolding::emitHelperTypes):
(WebCore::WHLSL::Metal::VertexEntryPointScaffolding::emitSignature):
(WebCore::WHLSL::Metal::VertexEntryPointScaffolding::emitUnpack):
(WebCore::WHLSL::Metal::VertexEntryPointScaffolding::emitPack):
(WebCore::WHLSL::Metal::FragmentEntryPointScaffolding::emitHelperTypes):
(WebCore::WHLSL::Metal::FragmentEntryPointScaffolding::emitSignature):
(WebCore::WHLSL::Metal::FragmentEntryPointScaffolding::emitUnpack):
(WebCore::WHLSL::Metal::FragmentEntryPointScaffolding::emitPack):
(WebCore::WHLSL::Metal::ComputeEntryPointScaffolding::emitHelperTypes):
(WebCore::WHLSL::Metal::ComputeEntryPointScaffolding::emitSignature):
(WebCore::WHLSL::Metal::ComputeEntryPointScaffolding::emitUnpack):
(WebCore::WHLSL::Metal::ComputeEntryPointScaffolding::emitPack):
* Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.h:
* Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.cpp:
(WebCore::WHLSL::Metal::FunctionDefinitionWriter::visit):
(WebCore::WHLSL::Metal::FunctionDefinitionWriter::emitLoop):
* Modules/webgpu/WHLSL/Metal/WHLSLNativeFunctionWriter.cpp:
(WebCore::WHLSL::Metal::inlineNativeFunction):
* Modules/webgpu/WHLSL/Metal/WHLSLNativeFunctionWriter.h:
* Modules/webgpu/WHLSL/Metal/WHLSLTypeNamer.cpp:
(WebCore::WHLSL::Metal::TypeNamer::emitUnnamedTypeDefinition):

Source/WTF:

Adds simple Indentation class to allow programatic indenting compatible with StringTypeAdapter
aware functions/class (e.g. makeString and StringBuilder). Templatized on the number of spaces
to indent.

Also adds IndentationScope, which uses RAII to increment/decrement the indentation value.

* wtf/text/StringConcatenate.h:
(WTF::Indentation::operator++):
(WTF::Indentation::operator--):
(WTF::IndentationScope::IndentationScope):
(WTF::IndentationScope::~IndentationScope):

git-svn-id: https://svn.webkit.org/repository/webkit/trunk@248892 268f45cc-cd09-0410-ab3c-d52691b4dbfc

Source/WTF/ChangeLog
Source/WTF/wtf/text/StringConcatenate.h
Source/WebCore/ChangeLog
Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp
Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.h
Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.cpp
Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLNativeFunctionWriter.cpp
Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLNativeFunctionWriter.h
Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLTypeNamer.cpp

index 7665951..3540454 100644 (file)
@@ -1,3 +1,22 @@
+2019-08-19  Sam Weinig  <weinig@apple.com>
+
+        [WHLSL] Make generated Metal code should be indented properly to ease reading while debugging
+        https://bugs.webkit.org/show_bug.cgi?id=200870
+
+        Reviewed by Saam Barati.
+
+        Adds simple Indentation class to allow programatic indenting compatible with StringTypeAdapter
+        aware functions/class (e.g. makeString and StringBuilder). Templatized on the number of spaces
+        to indent.
+        
+        Also adds IndentationScope, which uses RAII to increment/decrement the indentation value.
+
+        * wtf/text/StringConcatenate.h:
+        (WTF::Indentation::operator++):
+        (WTF::Indentation::operator--):
+        (WTF::IndentationScope::IndentationScope):
+        (WTF::IndentationScope::~IndentationScope):
+
 2019-08-19  Per Arne Vollan  <pvollan@apple.com>
 
         Unreviewed build fix for macOS 10.14 after r248832.
index 0537b10..dcf52db 100644 (file)
@@ -235,6 +235,58 @@ private:
     StringTypeAdapter<UnderlyingElementType> m_underlyingAdapter;
 };
 
+template<unsigned N>
+struct Indentation {
+    unsigned operator++() { return ++value; }
+    unsigned operator++(int) { return value++; }
+    unsigned operator--() { return --value; }
+    unsigned operator--(int) { return value--; }
+
+    unsigned value { 0 };
+};
+
+
+template<unsigned N>
+struct IndentationScope {
+    IndentationScope(Indentation<N>& indentation)
+        : m_indentation(indentation)
+    {
+        ++m_indentation;
+    }
+    ~IndentationScope()
+    {
+        --m_indentation;
+    }
+
+    Indentation<N>& m_indentation;
+};
+
+template<unsigned N> class StringTypeAdapter<Indentation<N>, void> {
+public:
+    StringTypeAdapter(Indentation<N> indentation)
+        : m_indentation { indentation }
+    {
+    }
+
+    unsigned length()
+    {
+        return m_indentation.value * N;
+    }
+
+    bool is8Bit()
+    {
+        return true;
+    }
+
+    template<typename CharacterType> void writeTo(CharacterType* destination)
+    {
+        std::fill_n(destination, m_indentation.value * N, ' ');
+    }
+
+private:
+    Indentation<N> m_indentation;
+};
+
 template<typename Adapter>
 inline bool are8Bit(Adapter adapter)
 {
@@ -308,6 +360,8 @@ String makeString(StringTypes... strings)
 
 } // namespace WTF
 
+using WTF::Indentation;
+using WTF::IndentationScope;
 using WTF::makeString;
 using WTF::pad;
 using WTF::tryMakeString;
index 9870c58..35a11b9 100644 (file)
@@ -1,3 +1,39 @@
+2019-08-19  Sam Weinig  <weinig@apple.com>
+
+        [WHLSL] Make generated Metal code should be indented properly to ease reading while debugging
+        https://bugs.webkit.org/show_bug.cgi?id=200870
+
+        Reviewed by Saam Barati.
+
+        Adds indentation of generated Metal code using the new WTF::Indentation and WTF::IndentationScope
+        classes. This makes reading the generated Metal code much easier while debugging it, and does not
+        seem to introduce any performance regression (I measured no difference on compute_boids.html test).
+
+        * Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp:
+        (WebCore::WHLSL::Metal::EntryPointScaffolding::emitResourceHelperTypes):
+        (WebCore::WHLSL::Metal::EntryPointScaffolding::emitUnpackResourcesAndNamedBuiltIns):
+        (WebCore::WHLSL::Metal::VertexEntryPointScaffolding::emitHelperTypes):
+        (WebCore::WHLSL::Metal::VertexEntryPointScaffolding::emitSignature):
+        (WebCore::WHLSL::Metal::VertexEntryPointScaffolding::emitUnpack):
+        (WebCore::WHLSL::Metal::VertexEntryPointScaffolding::emitPack):
+        (WebCore::WHLSL::Metal::FragmentEntryPointScaffolding::emitHelperTypes):
+        (WebCore::WHLSL::Metal::FragmentEntryPointScaffolding::emitSignature):
+        (WebCore::WHLSL::Metal::FragmentEntryPointScaffolding::emitUnpack):
+        (WebCore::WHLSL::Metal::FragmentEntryPointScaffolding::emitPack):
+        (WebCore::WHLSL::Metal::ComputeEntryPointScaffolding::emitHelperTypes):
+        (WebCore::WHLSL::Metal::ComputeEntryPointScaffolding::emitSignature):
+        (WebCore::WHLSL::Metal::ComputeEntryPointScaffolding::emitUnpack):
+        (WebCore::WHLSL::Metal::ComputeEntryPointScaffolding::emitPack):
+        * Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.h:
+        * Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.cpp:
+        (WebCore::WHLSL::Metal::FunctionDefinitionWriter::visit):
+        (WebCore::WHLSL::Metal::FunctionDefinitionWriter::emitLoop):
+        * Modules/webgpu/WHLSL/Metal/WHLSLNativeFunctionWriter.cpp:
+        (WebCore::WHLSL::Metal::inlineNativeFunction):
+        * Modules/webgpu/WHLSL/Metal/WHLSLNativeFunctionWriter.h:
+        * Modules/webgpu/WHLSL/Metal/WHLSLTypeNamer.cpp:
+        (WebCore::WHLSL::Metal::TypeNamer::emitUnnamedTypeDefinition):
+
 2019-08-19  Devin Rousso  <drousso@apple.com>
 
         Web Inspector: have more aggressive checks for dataURLs provided to `console.screenshot`
index 290daa2..7e75937 100644 (file)
@@ -139,39 +139,42 @@ EntryPointScaffolding::EntryPointScaffolding(AST::FunctionDefinition& functionDe
         m_parameterVariables.uncheckedAppend(m_generateNextVariableName());
 }
 
-void EntryPointScaffolding::emitResourceHelperTypes(StringBuilder& stringBuilder)
+void EntryPointScaffolding::emitResourceHelperTypes(StringBuilder& stringBuilder, Indentation<4> indent)
 {
     for (size_t i = 0; i < m_layout.size(); ++i) {
-        stringBuilder.append("struct ", m_namedBindGroups[i].structName, " {\n");
-        Vector<std::pair<unsigned, String>> structItems;
-        for (size_t j = 0; j < m_layout[i].bindings.size(); ++j) {
-            auto iterator = m_resourceMap.find(&m_layout[i].bindings[j]);
-            if (iterator == m_resourceMap.end())
-                continue;
-            auto& type = m_entryPointItems.inputs[iterator->value].unnamedType->unifyNode();
-            if (is<AST::UnnamedType>(type) && is<AST::ReferenceType>(downcast<AST::UnnamedType>(type))) {
-                auto& referenceType = downcast<AST::ReferenceType>(downcast<AST::UnnamedType>(type));
-                auto mangledTypeName = m_typeNamer.mangledNameForType(referenceType.elementType());
-                auto addressSpace = toString(referenceType.addressSpace());
-                auto elementName = m_namedBindGroups[i].namedBindings[j].elementName;
-                auto index = m_namedBindGroups[i].namedBindings[j].index;
-                structItems.append(std::make_pair(index, makeString(addressSpace, " ", mangledTypeName, "* ", elementName, " [[id(", index, ")]];")));
-                if (auto lengthInformation = m_namedBindGroups[i].namedBindings[j].lengthInformation)
-                    structItems.append(std::make_pair(lengthInformation->index, makeString("uint2 ", lengthInformation->elementName, " [[id(", lengthInformation->index, ")]];")));
-            } else if (is<AST::NamedType>(type) && is<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(type))) {
-                auto& namedType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(type));
-                auto mangledTypeName = m_typeNamer.mangledNameForType(namedType);
-                auto elementName = m_namedBindGroups[i].namedBindings[j].elementName;
-                auto index = m_namedBindGroups[i].namedBindings[j].index;
-                structItems.append(std::make_pair(index, makeString(mangledTypeName, ' ', elementName, " [[id(", index, ")]];")));
+        stringBuilder.append(indent, "struct ", m_namedBindGroups[i].structName, " {\n");
+        {
+            IndentationScope scope(indent);
+            Vector<std::pair<unsigned, String>> structItems;
+            for (size_t j = 0; j < m_layout[i].bindings.size(); ++j) {
+                auto iterator = m_resourceMap.find(&m_layout[i].bindings[j]);
+                if (iterator == m_resourceMap.end())
+                    continue;
+                auto& type = m_entryPointItems.inputs[iterator->value].unnamedType->unifyNode();
+                if (is<AST::UnnamedType>(type) && is<AST::ReferenceType>(downcast<AST::UnnamedType>(type))) {
+                    auto& referenceType = downcast<AST::ReferenceType>(downcast<AST::UnnamedType>(type));
+                    auto mangledTypeName = m_typeNamer.mangledNameForType(referenceType.elementType());
+                    auto addressSpace = toString(referenceType.addressSpace());
+                    auto elementName = m_namedBindGroups[i].namedBindings[j].elementName;
+                    auto index = m_namedBindGroups[i].namedBindings[j].index;
+                    structItems.append(std::make_pair(index, makeString(addressSpace, " ", mangledTypeName, "* ", elementName, " [[id(", index, ")]];")));
+                    if (auto lengthInformation = m_namedBindGroups[i].namedBindings[j].lengthInformation)
+                        structItems.append(std::make_pair(lengthInformation->index, makeString("uint2 ", lengthInformation->elementName, " [[id(", lengthInformation->index, ")]];")));
+                } else if (is<AST::NamedType>(type) && is<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(type))) {
+                    auto& namedType = downcast<AST::NativeTypeDeclaration>(downcast<AST::NamedType>(type));
+                    auto mangledTypeName = m_typeNamer.mangledNameForType(namedType);
+                    auto elementName = m_namedBindGroups[i].namedBindings[j].elementName;
+                    auto index = m_namedBindGroups[i].namedBindings[j].index;
+                    structItems.append(std::make_pair(index, makeString(mangledTypeName, ' ', elementName, " [[id(", index, ")]];")));
+                }
             }
+            std::sort(structItems.begin(), structItems.end(), [](const std::pair<unsigned, String>& left, const std::pair<unsigned, String>& right) {
+                return left.first < right.first;
+            });
+            for (const auto& structItem : structItems)
+                stringBuilder.append(indent, structItem.second, '\n');
         }
-        std::sort(structItems.begin(), structItems.end(), [](const std::pair<unsigned, String>& left, const std::pair<unsigned, String>& right) {
-            return left.first < right.first;
-        });
-        for (const auto& structItem : structItems)
-            stringBuilder.append("    ", structItem.second, '\n');
-        stringBuilder.append("};\n\n");
+        stringBuilder.append(indent, "};\n\n");
     }
 }
 
@@ -305,10 +308,10 @@ void EntryPointScaffolding::emitMangledOutputPath(StringBuilder& stringBuilder,
     }
 }
 
-void EntryPointScaffolding::emitUnpackResourcesAndNamedBuiltIns(StringBuilder& stringBuilder)
+void EntryPointScaffolding::emitUnpackResourcesAndNamedBuiltIns(StringBuilder& stringBuilder, Indentation<4> indent)
 {
     for (size_t i = 0; i < m_functionDefinition.parameters().size(); ++i)
-        stringBuilder.append(m_typeNamer.mangledNameForType(*m_functionDefinition.parameters()[i]->type()), ' ', m_parameterVariables[i], ";\n");
+        stringBuilder.append(indent, m_typeNamer.mangledNameForType(*m_functionDefinition.parameters()[i]->type()), ' ', m_parameterVariables[i], ";\n");
 
     for (size_t i = 0; i < m_layout.size(); ++i) {
         auto variableName = m_namedBindGroups[i].variableName;
@@ -326,13 +329,15 @@ void EntryPointScaffolding::emitUnpackResourcesAndNamedBuiltIns(StringBuilder& s
                 auto mangledTypeName = m_typeNamer.mangledNameForType(downcast<AST::ReferenceType>(unnamedType).elementType());
 
                 stringBuilder.append(
-                    "size_t ", lengthTemporaryName, " = ", variableName, '.', lengthElementName, ".y;\n",
-                    lengthTemporaryName, " = ", lengthTemporaryName, " << 32;\n",
-                    lengthTemporaryName, " = ", lengthTemporaryName, " | ", variableName, '.', lengthElementName, ".x;\n",
-                    lengthTemporaryName, " = ", lengthTemporaryName, " / sizeof(", mangledTypeName, ");\n",
-                    "if (", lengthTemporaryName, " > 0xFFFFFFFF) ", lengthTemporaryName, " = 0xFFFFFFFF;\n"
+                    indent, "size_t ", lengthTemporaryName, " = ", variableName, '.', lengthElementName, ".y;\n",
+                    indent, lengthTemporaryName, " = ", lengthTemporaryName, " << 32;\n",
+                    indent, lengthTemporaryName, " = ", lengthTemporaryName, " | ", variableName, '.', lengthElementName, ".x;\n",
+                    indent, lengthTemporaryName, " = ", lengthTemporaryName, " / sizeof(", mangledTypeName, ");\n",
+                    indent, "if (", lengthTemporaryName, " > 0xFFFFFFFF)\n",
+                    indent, "    ", lengthTemporaryName, " = 0xFFFFFFFF;\n"
                 );
-                    
+
+                stringBuilder.append(indent);
                 emitMangledInputPath(stringBuilder, path);
                 stringBuilder.append(
                     " = { ", variableName, '.', elementName, ", static_cast<uint32_t>(", lengthTemporaryName, ") };\n"
@@ -341,6 +346,7 @@ void EntryPointScaffolding::emitUnpackResourcesAndNamedBuiltIns(StringBuilder& s
                 auto& path = m_entryPointItems.inputs[iterator->value].path;
                 auto elementName = m_namedBindGroups[i].namedBindings[j].elementName;
                 
+                stringBuilder.append(indent);
                 emitMangledInputPath(stringBuilder, path);
                 stringBuilder.append(" = ", variableName, '.', elementName, ";\n");
             }
@@ -353,6 +359,7 @@ void EntryPointScaffolding::emitUnpackResourcesAndNamedBuiltIns(StringBuilder& s
         auto& variableName = namedBuiltIn.variableName;
         auto mangledTypeName = m_typeNamer.mangledNameForType(*item.unnamedType);
 
+        stringBuilder.append(indent);
         emitMangledInputPath(stringBuilder, path);
         stringBuilder.append(" = ", mangledTypeName, '(', variableName, ");\n");
     }
@@ -390,66 +397,72 @@ VertexEntryPointScaffolding::VertexEntryPointScaffolding(AST::FunctionDefinition
     }
 }
 
-void VertexEntryPointScaffolding::emitHelperTypes(StringBuilder& stringBuilder)
+void VertexEntryPointScaffolding::emitHelperTypes(StringBuilder& stringBuilder, Indentation<4> indent)
 {
-    stringBuilder.append("struct ", m_stageInStructName, " {\n");
-    for (auto& namedStageIn : m_namedStageIns) {
-        auto mangledTypeName = m_typeNamer.mangledNameForType(*m_entryPointItems.inputs[namedStageIn.indexInEntryPointItems].unnamedType);
-        auto elementName = namedStageIn.elementName;
-        auto attributeIndex = namedStageIn.attributeIndex;
-        stringBuilder.append("    ", mangledTypeName, ' ', elementName, " [[attribute(", attributeIndex, ")]];\n");
+    stringBuilder.append(indent, "struct ", m_stageInStructName, " {\n");
+    {
+        IndentationScope scope(indent);
+        for (auto& namedStageIn : m_namedStageIns) {
+            auto mangledTypeName = m_typeNamer.mangledNameForType(*m_entryPointItems.inputs[namedStageIn.indexInEntryPointItems].unnamedType);
+            auto elementName = namedStageIn.elementName;
+            auto attributeIndex = namedStageIn.attributeIndex;
+            stringBuilder.append(indent, mangledTypeName, ' ', elementName, " [[attribute(", attributeIndex, ")]];\n");
+        }
     }
     stringBuilder.append(
-        "};\n\n"
-        "struct ", m_returnStructName, " {\n"
+        indent, "};\n\n",
+        indent, "struct ", m_returnStructName, " {\n"
     );
-    for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
-        auto& outputItem = m_entryPointItems.outputs[i];
-        auto& internalTypeName = m_namedOutputs[i].internalTypeName;
-        auto elementName = m_namedOutputs[i].elementName;
-        auto attribute = attributeForSemantic(*outputItem.semantic);
-        stringBuilder.append("    ", internalTypeName, ' ', elementName, ' ', attribute, ";\n");
+    {
+        IndentationScope scope(indent);
+        for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
+            auto& outputItem = m_entryPointItems.outputs[i];
+            auto& internalTypeName = m_namedOutputs[i].internalTypeName;
+            auto elementName = m_namedOutputs[i].elementName;
+            auto attribute = attributeForSemantic(*outputItem.semantic);
+            stringBuilder.append(indent, internalTypeName, ' ', elementName, ' ', attribute, ";\n");
+        }
     }
-    stringBuilder.append(
-        "};\n\n"
-    );
+    stringBuilder.append(indent, "};\n\n");
     
-    emitResourceHelperTypes(stringBuilder);
+    emitResourceHelperTypes(stringBuilder, indent);
 }
 
-void VertexEntryPointScaffolding::emitSignature(StringBuilder& stringBuilder, MangledFunctionName functionName)
+void VertexEntryPointScaffolding::emitSignature(StringBuilder& stringBuilder, MangledFunctionName functionName, Indentation<4> indent)
 {
-    stringBuilder.append("vertex ", m_returnStructName, ' ', functionName, '(', m_stageInStructName, ' ', m_stageInParameterName, " [[stage_in]]");
+    stringBuilder.append(indent, "vertex ", m_returnStructName, ' ', functionName, '(', m_stageInStructName, ' ', m_stageInParameterName, " [[stage_in]]");
     emitResourceSignature(stringBuilder, IncludePrecedingComma::Yes);
     emitBuiltInsSignature(stringBuilder, IncludePrecedingComma::Yes);
-    stringBuilder.append(')');
+    stringBuilder.append(")\n");
 }
 
-void VertexEntryPointScaffolding::emitUnpack(StringBuilder& stringBuilder)
+void VertexEntryPointScaffolding::emitUnpack(StringBuilder& stringBuilder, Indentation<4> indent)
 {
-    emitUnpackResourcesAndNamedBuiltIns(stringBuilder);
+    emitUnpackResourcesAndNamedBuiltIns(stringBuilder, indent);
 
     for (auto& namedStageIn : m_namedStageIns) {
         auto& path = m_entryPointItems.inputs[namedStageIn.indexInEntryPointItems].path;
         auto& elementName = namedStageIn.elementName;
+        
+        stringBuilder.append(indent);
         emitMangledInputPath(stringBuilder, path);
         stringBuilder.append(" = ", m_stageInParameterName, '.', elementName, ";\n");
     }
 }
 
-void VertexEntryPointScaffolding::emitPack(StringBuilder& stringBuilder, MangledVariableName inputVariableName, MangledVariableName outputVariableName)
+void VertexEntryPointScaffolding::emitPack(StringBuilder& stringBuilder, MangledVariableName inputVariableName, MangledVariableName outputVariableName, Indentation<4> indent)
 {
-    stringBuilder.append(m_returnStructName, ' ', outputVariableName, ";\n");
+    stringBuilder.append(indent, m_returnStructName, ' ', outputVariableName, ";\n");
     if (m_entryPointItems.outputs.size() == 1 && !m_entryPointItems.outputs[0].path.size()) {
         auto& elementName = m_namedOutputs[0].elementName;
-        stringBuilder.append(outputVariableName, '.', elementName, " = ", inputVariableName, ";\n");
+        stringBuilder.append(indent, outputVariableName, '.', elementName, " = ", inputVariableName, ";\n");
         return;
     }
     for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
         auto& elementName = m_namedOutputs[i].elementName;
         auto& internalTypeName = m_namedOutputs[i].internalTypeName;
         auto& path = m_entryPointItems.outputs[i].path;
-        stringBuilder.append(outputVariableName, '.', elementName, " = ", internalTypeName, '(', inputVariableName);
+        stringBuilder.append(indent, outputVariableName, '.', elementName, " = ", internalTypeName, '(', inputVariableName);
         emitMangledOutputPath(stringBuilder, path);
         stringBuilder.append(");\n");
     }
@@ -489,66 +502,72 @@ FragmentEntryPointScaffolding::FragmentEntryPointScaffolding(AST::FunctionDefini
     }
 }
 
-void FragmentEntryPointScaffolding::emitHelperTypes(StringBuilder& stringBuilder)
+void FragmentEntryPointScaffolding::emitHelperTypes(StringBuilder& stringBuilder, Indentation<4> indent)
 {
-    stringBuilder.append("struct ", m_stageInStructName, " {\n");
-    for (auto& namedStageIn : m_namedStageIns) {
-        auto mangledTypeName = m_typeNamer.mangledNameForType(*m_entryPointItems.inputs[namedStageIn.indexInEntryPointItems].unnamedType);
-        auto elementName = namedStageIn.elementName;
-        auto attributeIndex = namedStageIn.attributeIndex;
-        stringBuilder.append("    ", mangledTypeName, ' ', elementName, " [[user(user", attributeIndex, ")]];\n");
+    stringBuilder.append(indent, "struct ", m_stageInStructName, " {\n");
+    {
+        IndentationScope scope(indent);
+        for (auto& namedStageIn : m_namedStageIns) {
+            auto mangledTypeName = m_typeNamer.mangledNameForType(*m_entryPointItems.inputs[namedStageIn.indexInEntryPointItems].unnamedType);
+            auto elementName = namedStageIn.elementName;
+            auto attributeIndex = namedStageIn.attributeIndex;
+            stringBuilder.append(indent, mangledTypeName, ' ', elementName, " [[user(user", attributeIndex, ")]];\n");
+        }
     }
     stringBuilder.append(
-        "};\n\n"
-        "struct ", m_returnStructName, " {\n"
+        indent, "};\n\n",
+        indent, "struct ", m_returnStructName, " {\n"
     );
-    for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
-        auto& outputItem = m_entryPointItems.outputs[i];
-        auto& internalTypeName = m_namedOutputs[i].internalTypeName;
-        auto elementName = m_namedOutputs[i].elementName;
-        auto attribute = attributeForSemantic(*outputItem.semantic);
-        stringBuilder.append("    ", internalTypeName, ' ', elementName, ' ', attribute, ";\n");
+    {
+        IndentationScope scope(indent);
+        for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
+            auto& outputItem = m_entryPointItems.outputs[i];
+            auto& internalTypeName = m_namedOutputs[i].internalTypeName;
+            auto elementName = m_namedOutputs[i].elementName;
+            auto attribute = attributeForSemantic(*outputItem.semantic);
+            stringBuilder.append(indent, internalTypeName, ' ', elementName, ' ', attribute, ";\n");
+        }
     }
-    stringBuilder.append(
-        "};\n\n"
-    );
+    stringBuilder.append(indent, "};\n\n");
 
-    emitResourceHelperTypes(stringBuilder);
+    emitResourceHelperTypes(stringBuilder, indent);
 }
 
-void FragmentEntryPointScaffolding::emitSignature(StringBuilder& stringBuilder, MangledFunctionName functionName)
+void FragmentEntryPointScaffolding::emitSignature(StringBuilder& stringBuilder, MangledFunctionName functionName, Indentation<4> indent)
 {
-    stringBuilder.append("fragment ", m_returnStructName, ' ', functionName, '(', m_stageInStructName, ' ', m_stageInParameterName, " [[stage_in]]");
+    stringBuilder.append(indent, "fragment ", m_returnStructName, ' ', functionName, '(', m_stageInStructName, ' ', m_stageInParameterName, " [[stage_in]]");
     emitResourceSignature(stringBuilder, IncludePrecedingComma::Yes);
     emitBuiltInsSignature(stringBuilder, IncludePrecedingComma::Yes);
-    stringBuilder.append(')');
+    stringBuilder.append(")\n");
 }
 
-void FragmentEntryPointScaffolding::emitUnpack(StringBuilder& stringBuilder)
+void FragmentEntryPointScaffolding::emitUnpack(StringBuilder& stringBuilder, Indentation<4> indent)
 {
-    emitUnpackResourcesAndNamedBuiltIns(stringBuilder);
+    emitUnpackResourcesAndNamedBuiltIns(stringBuilder, indent);
 
     for (auto& namedStageIn : m_namedStageIns) {
         auto& path = m_entryPointItems.inputs[namedStageIn.indexInEntryPointItems].path;
         auto& elementName = namedStageIn.elementName;
+
+        stringBuilder.append(indent);
         emitMangledInputPath(stringBuilder, path);
         stringBuilder.append(" = ", m_stageInParameterName, '.', elementName, ";\n");
     }
 }
 
-void FragmentEntryPointScaffolding::emitPack(StringBuilder& stringBuilder, MangledVariableName inputVariableName, MangledVariableName outputVariableName)
+void FragmentEntryPointScaffolding::emitPack(StringBuilder& stringBuilder, MangledVariableName inputVariableName, MangledVariableName outputVariableName, Indentation<4> indent)
 {
-    stringBuilder.append(m_returnStructName, ' ', outputVariableName, ";\n");
+    stringBuilder.append(indent, m_returnStructName, ' ', outputVariableName, ";\n");
     if (m_entryPointItems.outputs.size() == 1 && !m_entryPointItems.outputs[0].path.size()) {
         auto& elementName = m_namedOutputs[0].elementName;
-        stringBuilder.append(outputVariableName, '.', elementName, " = ", inputVariableName, ";\n");
+        stringBuilder.append(indent, outputVariableName, '.', elementName, " = ", inputVariableName, ";\n");
         return;
     }
     for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
         auto& elementName = m_namedOutputs[i].elementName;
         auto& internalTypeName = m_namedOutputs[i].internalTypeName;
         auto& path = m_entryPointItems.outputs[i].path;
-        stringBuilder.append(outputVariableName, '.', elementName, " = ", internalTypeName, '(', inputVariableName);
+        stringBuilder.append(indent, outputVariableName, '.', elementName, " = ", internalTypeName, '(', inputVariableName);
         emitMangledOutputPath(stringBuilder, path);
         stringBuilder.append(");\n");
     }
@@ -559,25 +578,25 @@ ComputeEntryPointScaffolding::ComputeEntryPointScaffolding(AST::FunctionDefiniti
 {
 }
 
-void ComputeEntryPointScaffolding::emitHelperTypes(StringBuilder& stringBuilder)
+void ComputeEntryPointScaffolding::emitHelperTypes(StringBuilder& stringBuilder, Indentation<4> indent)
 {
-    emitResourceHelperTypes(stringBuilder);
+    emitResourceHelperTypes(stringBuilder, indent);
 }
 
-void ComputeEntryPointScaffolding::emitSignature(StringBuilder& stringBuilder, MangledFunctionName functionName)
+void ComputeEntryPointScaffolding::emitSignature(StringBuilder& stringBuilder, MangledFunctionName functionName, Indentation<4> indent)
 {
-    stringBuilder.append("kernel void ", functionName, '(');
+    stringBuilder.append(indent, "kernel void ", functionName, '(');
     bool addedToSignature = emitResourceSignature(stringBuilder, IncludePrecedingComma::No);
     emitBuiltInsSignature(stringBuilder, addedToSignature ? IncludePrecedingComma::Yes : IncludePrecedingComma::No);
-    stringBuilder.append(')');
+    stringBuilder.append(")\n");
 }
 
-void ComputeEntryPointScaffolding::emitUnpack(StringBuilder& stringBuilder)
+void ComputeEntryPointScaffolding::emitUnpack(StringBuilder& stringBuilder, Indentation<4> indent)
 {
-    emitUnpackResourcesAndNamedBuiltIns(stringBuilder);
+    emitUnpackResourcesAndNamedBuiltIns(stringBuilder, indent);
 }
 
-void ComputeEntryPointScaffolding::emitPack(StringBuilder&, MangledVariableName, MangledVariableName)
+void ComputeEntryPointScaffolding::emitPack(StringBuilder&, MangledVariableName, MangledVariableName, Indentation<4>)
 {
     ASSERT_NOT_REACHED();
 }
index cbcc484..48a8c5b 100644 (file)
@@ -54,17 +54,17 @@ class EntryPointScaffolding {
 public:
     virtual ~EntryPointScaffolding() = default;
 
-    virtual void emitHelperTypes(StringBuilder&) = 0;
-    virtual void emitSignature(StringBuilder&, MangledFunctionName) = 0;
-    virtual void emitUnpack(StringBuilder&) = 0;
-    virtual void emitPack(StringBuilder&, MangledVariableName existingVariableName, MangledVariableName) = 0;
+    virtual void emitHelperTypes(StringBuilder&, Indentation<4>) = 0;
+    virtual void emitSignature(StringBuilder&, MangledFunctionName, Indentation<4>) = 0;
+    virtual void emitUnpack(StringBuilder&, Indentation<4>) = 0;
+    virtual void emitPack(StringBuilder&, MangledVariableName existingVariableName, MangledVariableName, Indentation<4>) = 0;
 
     Vector<MangledVariableName>& parameterVariables() { return m_parameterVariables; }
 
 protected:
     EntryPointScaffolding(AST::FunctionDefinition&, Intrinsics&, TypeNamer&, EntryPointItems&, HashMap<Binding*, size_t>& resourceMap, Layout&, std::function<MangledVariableName()>&& generateNextVariableName);
 
-    void emitResourceHelperTypes(StringBuilder&);
+    void emitResourceHelperTypes(StringBuilder&, Indentation<4>);
 
     enum class IncludePrecedingComma {
         Yes,
@@ -75,7 +75,7 @@ protected:
 
     void emitMangledInputPath(StringBuilder&, Vector<String>& path);
     void emitMangledOutputPath(StringBuilder&, Vector<String>& path);
-    void emitUnpackResourcesAndNamedBuiltIns(StringBuilder&);
+    void emitUnpackResourcesAndNamedBuiltIns(StringBuilder&, Indentation<4>);
 
     AST::FunctionDefinition& m_functionDefinition;
     Intrinsics& m_intrinsics;
@@ -118,10 +118,10 @@ public:
     virtual ~VertexEntryPointScaffolding() = default;
 
 private:
-    void emitHelperTypes(StringBuilder&) override;
-    void emitSignature(StringBuilder&, MangledFunctionName) override;
-    void emitUnpack(StringBuilder&) override;
-    void emitPack(StringBuilder&, MangledVariableName existingVariableName, MangledVariableName) override;
+    void emitHelperTypes(StringBuilder&, Indentation<4>) override;
+    void emitSignature(StringBuilder&, MangledFunctionName, Indentation<4>) override;
+    void emitUnpack(StringBuilder&, Indentation<4>) override;
+    void emitPack(StringBuilder&, MangledVariableName existingVariableName, MangledVariableName, Indentation<4>) override;
 
     HashMap<VertexAttribute*, size_t>& m_matchedVertexAttributes;
     MangledTypeName m_stageInStructName;
@@ -148,10 +148,10 @@ public:
     virtual ~FragmentEntryPointScaffolding() = default;
 
 private:
-    void emitHelperTypes(StringBuilder&) override;
-    void emitSignature(StringBuilder&, MangledFunctionName) override;
-    void emitUnpack(StringBuilder&) override;
-    void emitPack(StringBuilder&, MangledVariableName existingVariableName, MangledVariableName) override;
+    void emitHelperTypes(StringBuilder&, Indentation<4>) override;
+    void emitSignature(StringBuilder&, MangledFunctionName, Indentation<4>) override;
+    void emitUnpack(StringBuilder&, Indentation<4>) override;
+    void emitPack(StringBuilder&, MangledVariableName existingVariableName, MangledVariableName, Indentation<4>) override;
 
     MangledTypeName m_stageInStructName;
     MangledTypeName m_returnStructName;
@@ -177,10 +177,10 @@ public:
     virtual ~ComputeEntryPointScaffolding() = default;
 
 private:
-    void emitHelperTypes(StringBuilder&) override;
-    void emitSignature(StringBuilder&, MangledFunctionName) override;
-    void emitUnpack(StringBuilder&) override;
-    void emitPack(StringBuilder&, MangledVariableName existingVariableName, MangledVariableName) override;
+    void emitHelperTypes(StringBuilder&, Indentation<4>) override;
+    void emitSignature(StringBuilder&, MangledFunctionName, Indentation<4>) override;
+    void emitUnpack(StringBuilder&, Indentation<4>) override;
+    void emitPack(StringBuilder&, MangledVariableName existingVariableName, MangledVariableName, Indentation<4>) override;
 };
 
 }
index 5dd9c8e..85ab8d9 100644 (file)
@@ -205,6 +205,7 @@ protected:
     Layout& m_layout;
     unsigned m_variableCount { 0 };
     Optional<MangledVariableName> m_breakOutOfCurrentLoopEarlyVariable;
+    Indentation<4> m_indent { 0 };
 };
 
 void FunctionDefinitionWriter::visit(AST::NativeFunctionDeclaration&)
@@ -221,23 +222,28 @@ void FunctionDefinitionWriter::visit(AST::FunctionDefinition& functionDefinition
         if (!entryPointScaffolding)
             return;
         m_entryPointScaffolding = WTFMove(entryPointScaffolding);
-        m_entryPointScaffolding->emitHelperTypes(m_stringBuilder);
-        m_stringBuilder.append('\n');
-        m_entryPointScaffolding->emitSignature(m_stringBuilder, iterator->value);
-        m_stringBuilder.append(" {\n");
-        m_entryPointScaffolding->emitUnpack(m_stringBuilder);
-    
-        for (size_t i = 0; i < functionDefinition.parameters().size(); ++i) {
-            auto addResult = m_variableMapping.add(&functionDefinition.parameters()[i], m_entryPointScaffolding->parameterVariables()[i]);
-            ASSERT_UNUSED(addResult, addResult.isNewEntry);
+        
+        m_entryPointScaffolding->emitHelperTypes(m_stringBuilder, m_indent);
+        m_entryPointScaffolding->emitSignature(m_stringBuilder, iterator->value, m_indent);
+        m_stringBuilder.append(m_indent, "{\n");
+        {
+            IndentationScope scope(m_indent);
+
+            m_entryPointScaffolding->emitUnpack(m_stringBuilder, m_indent);
+        
+            for (size_t i = 0; i < functionDefinition.parameters().size(); ++i) {
+                auto addResult = m_variableMapping.add(&functionDefinition.parameters()[i], m_entryPointScaffolding->parameterVariables()[i]);
+                ASSERT_UNUSED(addResult, addResult.isNewEntry);
+            }
+            checkErrorAndVisit(functionDefinition.block());
+            ASSERT(m_stack.isEmpty());
         }
-        checkErrorAndVisit(functionDefinition.block());
-        ASSERT(m_stack.isEmpty());
-        m_stringBuilder.append("}\n");
+        m_stringBuilder.append("}\n\n");
+
         m_entryPointScaffolding = nullptr;
     } else {
         ASSERT(m_entryPointScaffolding == nullptr);
-        m_stringBuilder.append(m_typeNamer.mangledNameForType(functionDefinition.type()), ' ', iterator->value, '(');
+        m_stringBuilder.append(m_indent, m_typeNamer.mangledNameForType(functionDefinition.type()), ' ', iterator->value, '(');
         for (size_t i = 0; i < functionDefinition.parameters().size(); ++i) {
             auto& parameter = functionDefinition.parameters()[i];
             if (i)
@@ -266,10 +272,13 @@ void FunctionDefinitionWriter::visit(AST::Statement& statement)
 
 void FunctionDefinitionWriter::visit(AST::Block& block)
 {
-    m_stringBuilder.append("{\n");
-    for (auto& statement : block.statements())
-        checkErrorAndVisit(statement);
-    m_stringBuilder.append("}\n");
+    m_stringBuilder.append(m_indent, "{\n");
+    {
+        IndentationScope scope(m_indent);
+        for (auto& statement : block.statements())
+            checkErrorAndVisit(statement);
+    }
+    m_stringBuilder.append(m_indent, "}\n");
 }
 
 void FunctionDefinitionWriter::visit(AST::Break&)
@@ -277,13 +286,13 @@ void FunctionDefinitionWriter::visit(AST::Break&)
     ASSERT(m_currentBreakContext);
     switch (*m_currentBreakContext) {
     case BreakContext::Switch:
-        m_stringBuilder.append("break;\n");
+        m_stringBuilder.append(m_indent, "break;\n");
         break;
     case BreakContext::Loop:
         ASSERT(m_breakOutOfCurrentLoopEarlyVariable);
         m_stringBuilder.append(
-            *m_breakOutOfCurrentLoopEarlyVariable, " = true;\n"
-            "break;\n"
+            m_indent, *m_breakOutOfCurrentLoopEarlyVariable, " = true;\n",
+            m_indent, "break;\n"
         );
         break;
     }
@@ -292,7 +301,7 @@ void FunctionDefinitionWriter::visit(AST::Break&)
 void FunctionDefinitionWriter::visit(AST::Continue&)
 {
     ASSERT(m_breakOutOfCurrentLoopEarlyVariable);
-    m_stringBuilder.append("break;\n");
+    m_stringBuilder.append(m_indent, "break;\n");
 }
 
 void FunctionDefinitionWriter::visit(AST::EffectfulExpressionStatement& effectfulExpressionStatement)
@@ -303,7 +312,7 @@ void FunctionDefinitionWriter::visit(AST::EffectfulExpressionStatement& effectfu
 
 void FunctionDefinitionWriter::visit(AST::Fallthrough&)
 {
-    m_stringBuilder.append("[[clang::fallthrough]];\n"); // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195808 Make sure this is okay. Alternatively, we could do nothing and just return here instead.
+    m_stringBuilder.append(m_indent, "[[clang::fallthrough]];\n"); // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195808 Make sure this is okay. Alternatively, we could do nothing and just return here instead.
 }
 
 void FunctionDefinitionWriter::emitLoop(LoopConditionLocation loopConditionLocation, AST::Expression* conditionExpression, AST::Expression* increment, AST::Statement& body)
@@ -311,36 +320,48 @@ void FunctionDefinitionWriter::emitLoop(LoopConditionLocation loopConditionLocat
     SetForScope<Optional<MangledVariableName>> loopVariableScope(m_breakOutOfCurrentLoopEarlyVariable, generateNextVariableName());
 
     m_stringBuilder.append(
-        "bool ", *m_breakOutOfCurrentLoopEarlyVariable, " = false;\n",
-        "while (true) {\n"
+        m_indent, "bool ", *m_breakOutOfCurrentLoopEarlyVariable, " = false;\n",
+        m_indent, "while (true) {\n"
     );
+    {
+        IndentationScope whileScope(m_indent);
 
-    if (loopConditionLocation == LoopConditionLocation::BeforeBody && conditionExpression) {
-        checkErrorAndVisit(*conditionExpression);
-        m_stringBuilder.append("if (!", takeLastValue(), ") break;\n");
-    }
+        if (loopConditionLocation == LoopConditionLocation::BeforeBody && conditionExpression) {
+            checkErrorAndVisit(*conditionExpression);
+            m_stringBuilder.append(
+                m_indent, "if (!", takeLastValue(), ")\n",
+                m_indent, "    break;\n");
+        }
 
-    m_stringBuilder.append("do {\n");
-    SetForScope<Optional<BreakContext>> breakContext(m_currentBreakContext, BreakContext::Loop);
-    checkErrorAndVisit(body);
-    m_stringBuilder.append(
-        "} while(false); \n"
-        "if (", *m_breakOutOfCurrentLoopEarlyVariable, ") break;\n"
-    );
+        m_stringBuilder.append(m_indent, "do {\n");
+        SetForScope<Optional<BreakContext>> breakContext(m_currentBreakContext, BreakContext::Loop);
 
-    if (increment) {
-        checkErrorAndVisit(*increment);
-        // Expression results get pushed to m_stack. We don't use the result
-        // of increment, so we dispense of that now.
-        takeLastValue();
-    }
+        {
+            IndentationScope doScope(m_indent);
+            checkErrorAndVisit(body);
+        }
+        m_stringBuilder.append(m_indent, "} while(false); \n");
 
-    if (loopConditionLocation == LoopConditionLocation::AfterBody && conditionExpression) {
-        checkErrorAndVisit(*conditionExpression);
-        m_stringBuilder.append("if (!", takeLastValue(), ") break;\n");
+        m_stringBuilder.append(
+            m_indent, "if (", *m_breakOutOfCurrentLoopEarlyVariable, ")\n",
+            m_indent, "    break;\n");
+
+        if (increment) {
+            checkErrorAndVisit(*increment);
+            // Expression results get pushed to m_stack. We don't use the result
+            // of increment, so we dispense of that now.
+            takeLastValue();
+        }
+
+        if (loopConditionLocation == LoopConditionLocation::AfterBody && conditionExpression) {
+            checkErrorAndVisit(*conditionExpression);
+            m_stringBuilder.append(
+                m_indent, "if (!", takeLastValue(), ")\n",
+                m_indent, "    break;\n");
+        }
     }
 
-    m_stringBuilder.append("} \n");
+    m_stringBuilder.append(m_indent, "} \n");
 }
 
 void FunctionDefinitionWriter::visit(AST::DoWhileLoop& doWhileLoop)
@@ -355,56 +376,69 @@ void FunctionDefinitionWriter::visit(AST::WhileLoop& whileLoop)
 
 void FunctionDefinitionWriter::visit(AST::ForLoop& forLoop)
 {
-    m_stringBuilder.append("{\n");
-    checkErrorAndVisit(forLoop.initialization());
-    emitLoop(LoopConditionLocation::BeforeBody, forLoop.condition(), forLoop.increment(), forLoop.body());
-    m_stringBuilder.append("}\n");
+    m_stringBuilder.append(m_indent, "{\n");
+    {
+        IndentationScope scope(m_indent);
+        checkErrorAndVisit(forLoop.initialization());
+        emitLoop(LoopConditionLocation::BeforeBody, forLoop.condition(), forLoop.increment(), forLoop.body());
+    }
+    m_stringBuilder.append(m_indent, "}\n");
 }
 
 void FunctionDefinitionWriter::visit(AST::IfStatement& ifStatement)
 {
     checkErrorAndVisit(ifStatement.conditional());
-    m_stringBuilder.append("if (", takeLastValue(), ") {\n");
-    checkErrorAndVisit(ifStatement.body());
+    m_stringBuilder.append(m_indent, "if (", takeLastValue(), ") {\n");
+    {
+        IndentationScope ifScope(m_indent);
+        checkErrorAndVisit(ifStatement.body());
+    }
     if (ifStatement.elseBody()) {
-        m_stringBuilder.append("} else {\n");
-        checkErrorAndVisit(*ifStatement.elseBody());
+        m_stringBuilder.append(m_indent, "} else {\n");
+        {
+            IndentationScope elseScope(m_indent);
+            checkErrorAndVisit(*ifStatement.elseBody());
+        }
     }
-    m_stringBuilder.append("}\n");
+    m_stringBuilder.append(m_indent, "}\n");
 }
 
 void FunctionDefinitionWriter::visit(AST::Return& returnStatement)
 {
     if (returnStatement.value()) {
         checkErrorAndVisit(*returnStatement.value());
+
         if (m_entryPointScaffolding) {
             auto variableName = generateNextVariableName();
-            m_entryPointScaffolding->emitPack(m_stringBuilder, takeLastValue(), variableName);
-            m_stringBuilder.append("return ", variableName, ";\n");
+            m_entryPointScaffolding->emitPack(m_stringBuilder, takeLastValue(), variableName, m_indent);
+            m_stringBuilder.append(m_indent, "return ", variableName, ";\n");
         } else
-            m_stringBuilder.append("return ", takeLastValue(), ";\n");
+            m_stringBuilder.append(m_indent, "return ", takeLastValue(), ";\n");
     } else
-        m_stringBuilder.append("return;\n");
+        m_stringBuilder.append(m_indent, "return;\n");
 }
 
 void FunctionDefinitionWriter::visit(AST::SwitchStatement& switchStatement)
 {
     checkErrorAndVisit(switchStatement.value());
 
-    m_stringBuilder.append("switch (", takeLastValue(), ") {");
-    for (auto& switchCase : switchStatement.switchCases())
-        checkErrorAndVisit(switchCase);
-    m_stringBuilder.append("}\n");
+    m_stringBuilder.append(m_indent, "switch (", takeLastValue(), ") {");
+    {
+        IndentationScope switchScope(m_indent);
+        for (auto& switchCase : switchStatement.switchCases())
+            checkErrorAndVisit(switchCase);
+    }
+    m_stringBuilder.append(m_indent, "}\n");
 }
 
 void FunctionDefinitionWriter::visit(AST::SwitchCase& switchCase)
 {
     if (switchCase.value()) {
-        m_stringBuilder.append("case ");
+        m_stringBuilder.append(m_indent, "case ");
         emitConstantExpressionString(*switchCase.value());
         m_stringBuilder.append(":\n");
     } else
-        m_stringBuilder.append("default:\n");
+        m_stringBuilder.append(m_indent, "default:\n");
     SetForScope<Optional<BreakContext>> breakContext(m_currentBreakContext, BreakContext::Switch);
     checkErrorAndVisit(switchCase.block());
     // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195812 Figure out whether we need to break or fallthrough.
@@ -419,7 +453,7 @@ void FunctionDefinitionWriter::visit(AST::IntegerLiteral& integerLiteral)
 {
     auto variableName = generateNextVariableName();
     auto mangledTypeName = m_typeNamer.mangledNameForType(integerLiteral.resolvedType());
-    m_stringBuilder.append(mangledTypeName, ' ', variableName, " = static_cast<", mangledTypeName, ">(", integerLiteral.value(), ");\n");
+    m_stringBuilder.append(m_indent, mangledTypeName, ' ', variableName, " = static_cast<", mangledTypeName, ">(", integerLiteral.value(), ");\n");
     appendRightValue(integerLiteral, variableName);
 }
 
@@ -427,7 +461,7 @@ void FunctionDefinitionWriter::visit(AST::UnsignedIntegerLiteral& unsignedIntege
 {
     auto variableName = generateNextVariableName();
     auto mangledTypeName = m_typeNamer.mangledNameForType(unsignedIntegerLiteral.resolvedType());
-    m_stringBuilder.append(mangledTypeName, ' ', variableName, " = static_cast<", mangledTypeName, ">(", unsignedIntegerLiteral.value(), ");\n");
+    m_stringBuilder.append(m_indent, mangledTypeName, ' ', variableName, " = static_cast<", mangledTypeName, ">(", unsignedIntegerLiteral.value(), ");\n");
     appendRightValue(unsignedIntegerLiteral, variableName);
 }
 
@@ -435,7 +469,7 @@ void FunctionDefinitionWriter::visit(AST::FloatLiteral& floatLiteral)
 {
     auto variableName = generateNextVariableName();
     auto mangledTypeName = m_typeNamer.mangledNameForType(floatLiteral.resolvedType());
-    m_stringBuilder.append(mangledTypeName, ' ', variableName, " = static_cast<", mangledTypeName, ">(", floatLiteral.value(), ");\n");
+    m_stringBuilder.append(m_indent, mangledTypeName, ' ', variableName, " = static_cast<", mangledTypeName, ">(", floatLiteral.value(), ");\n");
     appendRightValue(floatLiteral, variableName);
 }
 
@@ -446,7 +480,7 @@ void FunctionDefinitionWriter::visit(AST::NullLiteral& nullLiteral)
     bool isArrayReferenceType = is<AST::ArrayReferenceType>(unnamedType);
 
     auto variableName = generateNextVariableName();
-    m_stringBuilder.append(m_typeNamer.mangledNameForType(nullLiteral.resolvedType()), ' ', variableName, " = ");
+    m_stringBuilder.append(m_indent, m_typeNamer.mangledNameForType(nullLiteral.resolvedType()), ' ', variableName, " = ");
     if (isArrayReferenceType)
         m_stringBuilder.append("{ nullptr, 0 };\n");
     else
@@ -458,7 +492,7 @@ void FunctionDefinitionWriter::visit(AST::BooleanLiteral& booleanLiteral)
 {
     auto variableName = generateNextVariableName();
     auto mangledTypeName = m_typeNamer.mangledNameForType(booleanLiteral.resolvedType());
-    m_stringBuilder.append(mangledTypeName, ' ', variableName, " = static_cast<", mangledTypeName, ">(", booleanLiteral.value() ? "true" : "false", ");\n");
+    m_stringBuilder.append(m_indent, mangledTypeName, ' ', variableName, " = static_cast<", mangledTypeName, ">(", booleanLiteral.value() ? "true" : "false", ");\n");
     appendRightValue(booleanLiteral, variableName);
 }
 
@@ -468,7 +502,7 @@ void FunctionDefinitionWriter::visit(AST::EnumerationMemberLiteral& enumerationM
     ASSERT(enumerationMemberLiteral.enumerationDefinition());
     auto variableName = generateNextVariableName();
     auto mangledTypeName = m_typeNamer.mangledNameForType(enumerationMemberLiteral.resolvedType());
-    m_stringBuilder.append(mangledTypeName, ' ', variableName, " = ", mangledTypeName, "::", m_typeNamer.mangledNameForEnumerationMember(*enumerationMemberLiteral.enumerationMember()), ";\n");
+    m_stringBuilder.append(m_indent, mangledTypeName, ' ', variableName, " = ", mangledTypeName, "::", m_typeNamer.mangledNameForEnumerationMember(*enumerationMemberLiteral.enumerationMember()), ";\n");
     appendRightValue(enumerationMemberLiteral, variableName);
 }
 
@@ -492,8 +526,8 @@ void FunctionDefinitionWriter::visit(AST::GlobalVariableReference& globalVariabl
     auto mangledTypeName = m_typeNamer.mangledNameForType(globalVariableReference.resolvedType());
     checkErrorAndVisit(globalVariableReference.base());
     m_stringBuilder.append(
-        "thread ", mangledTypeName, "* ", pointerName, " = &", takeLastValue(), "->", m_typeNamer.mangledNameForStructureElement(globalVariableReference.structField()), ";\n",
-        mangledTypeName, ' ', valueName, " = ", "*", pointerName, ";\n"
+        m_indent, "thread ", mangledTypeName, "* ", pointerName, " = &", takeLastValue(), "->", m_typeNamer.mangledNameForStructureElement(globalVariableReference.structField()), ";\n",
+        m_indent, mangledTypeName, ' ', valueName, " = ", "*", pointerName, ";\n"
     );
     appendLeftValue(globalVariableReference, valueName, pointerName, Nullability::NotNull);
 }
@@ -523,9 +557,9 @@ void FunctionDefinitionWriter::visit(AST::VariableDeclaration& variableDeclarati
     // FIXME: https://bugs.webkit.org/show_bug.cgi?id=198160 Implement qualifiers.
     if (variableDeclaration.initializer()) {
         checkErrorAndVisit(*variableDeclaration.initializer());
-        m_stringBuilder.append(m_typeNamer.mangledNameForType(*variableDeclaration.type()), ' ', variableName, " = ", takeLastValue(), ";\n");
+        m_stringBuilder.append(m_indent, m_typeNamer.mangledNameForType(*variableDeclaration.type()), ' ', variableName, " = ", takeLastValue(), ";\n");
     } else
-        m_stringBuilder.append(m_typeNamer.mangledNameForType(*variableDeclaration.type()), ' ', variableName, " = { };\n");
+        m_stringBuilder.append(m_indent, m_typeNamer.mangledNameForType(*variableDeclaration.type()), ' ', variableName, " = { };\n");
 }
 
 void FunctionDefinitionWriter::visit(AST::AssignmentExpression& assignmentExpression)
@@ -534,10 +568,13 @@ void FunctionDefinitionWriter::visit(AST::AssignmentExpression& assignmentExpres
     auto [pointerName, nullability] = takeLastLeftValue();
     checkErrorAndVisit(assignmentExpression.right());
     auto [rightName, rightNullability] = takeLastValueAndNullability();
+
     if (nullability == Nullability::CanBeNull)
-        m_stringBuilder.append("if (", pointerName, ") *", pointerName, " = ", rightName, ";\n");
+        m_stringBuilder.append(
+            m_indent, "if (", pointerName, ")\n",
+            m_indent, "    *", pointerName, " = ", rightName, ";\n");
     else
-        m_stringBuilder.append("*", pointerName, " = ", rightName, ";\n");
+        m_stringBuilder.append(m_indent, "*", pointerName, " = ", rightName, ";\n");
     appendRightValueWithNullability(assignmentExpression, rightName, rightNullability);
 }
 
@@ -553,12 +590,20 @@ void FunctionDefinitionWriter::visit(AST::CallExpression& callExpression)
     MangledVariableName returnName;
     if (!isVoid) {
         returnName = generateNextVariableName();
-        m_stringBuilder.append(m_typeNamer.mangledNameForType(callExpression.resolvedType()), ' ', returnName, ";\n");
+        m_stringBuilder.append(m_indent, m_typeNamer.mangledNameForType(callExpression.resolvedType()), ' ', returnName, ";\n");
     }
 
-    if (is<AST::NativeFunctionDeclaration>(callExpression.function()))
-        inlineNativeFunction(m_stringBuilder, downcast<AST::NativeFunctionDeclaration>(callExpression.function()), returnName, argumentNames, m_intrinsics, m_typeNamer);
-    else {
+    if (is<AST::NativeFunctionDeclaration>(callExpression.function())) {
+        auto generateNextVariableName = [this]() -> MangledVariableName {
+            return this->generateNextVariableName();
+        };
+
+        m_stringBuilder.append('\n');
+        inlineNativeFunction(m_stringBuilder, downcast<AST::NativeFunctionDeclaration>(callExpression.function()), returnName, argumentNames, m_intrinsics, m_typeNamer, WTFMove(generateNextVariableName), m_indent);
+        m_stringBuilder.append('\n');
+    } else {
+        m_stringBuilder.append(m_indent);
+
         auto iterator = m_functionMapping.find(&callExpression.function());
         ASSERT(iterator != m_functionMapping.end());
         if (!isVoid)
@@ -592,15 +637,19 @@ void FunctionDefinitionWriter::visit(AST::DereferenceExpression& dereferenceExpr
     auto [inputPointer, nullability] = takeLastValueAndNullability();
     auto resultValue = generateNextVariableName();
     auto resultPointer = generateNextVariableName();
+
     m_stringBuilder.append(
-        m_typeNamer.mangledNameForType(dereferenceExpression.pointer().resolvedType()), ' ', resultPointer, " = ", inputPointer, ";\n",
-        m_typeNamer.mangledNameForType(dereferenceExpression.resolvedType()), ' ', resultValue, ";\n");
+        m_indent, m_typeNamer.mangledNameForType(dereferenceExpression.pointer().resolvedType()), ' ', resultPointer, " = ", inputPointer, ";\n",
+        m_indent, m_typeNamer.mangledNameForType(dereferenceExpression.resolvedType()), ' ', resultValue, ";\n");
     if (nullability == Nullability::CanBeNull) {
         m_stringBuilder.append(
-            "if (", resultPointer, ") ", resultValue, " = *", inputPointer, ";\n",
-            "else ", resultValue, " = { };\n");
+            m_indent, "if (", resultPointer, ")\n",
+            m_indent, "    ", resultValue, " = *", inputPointer, ";\n",
+            m_indent, "else\n",
+            m_indent, "    ", resultValue, " = { };\n"
+        );
     } else
-        m_stringBuilder.append(resultValue, " = *", inputPointer, ";\n");
+        m_stringBuilder.append(m_indent, resultValue, " = *", inputPointer, ";\n");
     appendLeftValue(dereferenceExpression, resultValue, resultPointer, nullability);
 }
 
@@ -611,7 +660,9 @@ void FunctionDefinitionWriter::visit(AST::LogicalExpression& logicalExpression)
     checkErrorAndVisit(logicalExpression.right());
     auto right = takeLastValue();
     auto variableName = generateNextVariableName();
-    m_stringBuilder.append(m_typeNamer.mangledNameForType(logicalExpression.resolvedType()), ' ', variableName, " = ", left);
+
+    m_stringBuilder.append(
+        m_indent, m_typeNamer.mangledNameForType(logicalExpression.resolvedType()), ' ', variableName, " = ", left);
     switch (logicalExpression.type()) {
     case AST::LogicalExpression::Type::And:
         m_stringBuilder.append(" && ");
@@ -630,7 +681,9 @@ void FunctionDefinitionWriter::visit(AST::LogicalNotExpression& logicalNotExpres
     checkErrorAndVisit(logicalNotExpression.operand());
     auto operand = takeLastValue();
     auto variableName = generateNextVariableName();
-    m_stringBuilder.append(m_typeNamer.mangledNameForType(logicalNotExpression.resolvedType()), ' ', variableName, " = !", operand, ";\n");
+
+    m_stringBuilder.append(
+        m_indent, m_typeNamer.mangledNameForType(logicalNotExpression.resolvedType()), ' ', variableName, " = !", operand, ";\n");
     appendRightValue(logicalNotExpression, variableName);
 }
 
@@ -645,17 +698,19 @@ void FunctionDefinitionWriter::visit(AST::MakeArrayReferenceExpression& makeArra
     if (is<AST::PointerType>(makeArrayReferenceExpression.leftValue().resolvedType())) {
         auto ptrValue = takeLastValue();
         m_stringBuilder.append(
-            mangledTypeName, ' ', variableName, ";\n",
-            "if (", ptrValue, ") ", variableName, " = { ", ptrValue, ", 1};\n",
-            "else ", variableName, " = { nullptr, 0 };\n"
+            m_indent, mangledTypeName, ' ', variableName, ";\n",
+            m_indent, "if (", ptrValue, ")\n",
+            m_indent, "    ", variableName, " = { ", ptrValue, ", 1};\n",
+            m_indent, "else\n",
+            m_indent, "    ", variableName, " = { nullptr, 0 };\n"
         );
     } else if (is<AST::ArrayType>(makeArrayReferenceExpression.leftValue().resolvedType())) {
         auto lValue = takeLastLeftValue().value;
         auto& arrayType = downcast<AST::ArrayType>(makeArrayReferenceExpression.leftValue().resolvedType());
-        m_stringBuilder.append(mangledTypeName, ' ', variableName, " = { ", lValue, "->data(), ", arrayType.numElements(), " };\n");
+        m_stringBuilder.append(m_indent, mangledTypeName, ' ', variableName, " = { ", lValue, "->data(), ", arrayType.numElements(), " };\n");
     } else {
         auto lValue = takeLastLeftValue().value;
-        m_stringBuilder.append(mangledTypeName, ' ', variableName, " = { ", lValue, ", 1 };\n");
+        m_stringBuilder.append(m_indent, mangledTypeName, ' ', variableName, " = { ", lValue, ", 1 };\n");
     }
     appendRightValue(makeArrayReferenceExpression, variableName);
 }
@@ -665,7 +720,7 @@ void FunctionDefinitionWriter::visit(AST::MakePointerExpression& makePointerExpr
     checkErrorAndVisit(makePointerExpression.leftValue());
     auto [pointer, nullability] = takeLastLeftValue();
     auto variableName = generateNextVariableName();
-    m_stringBuilder.append(m_typeNamer.mangledNameForType(makePointerExpression.resolvedType()), ' ', variableName, " = ", pointer, ";\n");
+    m_stringBuilder.append(m_indent, m_typeNamer.mangledNameForType(makePointerExpression.resolvedType()), ' ', variableName, " = ", pointer, ";\n");
     appendRightValueWithNullability(makePointerExpression, variableName, nullability);
 }
 
@@ -685,7 +740,7 @@ void FunctionDefinitionWriter::visit(AST::TernaryExpression& ternaryExpression)
     auto elseBody = takeLastValue();
 
     auto variableName = generateNextVariableName();
-    m_stringBuilder.append(m_typeNamer.mangledNameForType(ternaryExpression.resolvedType()), ' ', variableName, " = ", check, " ? ", body, " : ", elseBody, ";\n");
+    m_stringBuilder.append(m_indent, m_typeNamer.mangledNameForType(ternaryExpression.resolvedType()), ' ', variableName, " = ", check, " ? ", body, " : ", elseBody, ";\n");
     appendRightValue(ternaryExpression, variableName);
 }
 
@@ -694,8 +749,9 @@ void FunctionDefinitionWriter::visit(AST::VariableReference& variableReference)
     ASSERT(variableReference.variable());
     auto iterator = m_variableMapping.find(variableReference.variable());
     ASSERT(iterator != m_variableMapping.end());
+
     auto pointerName = generateNextVariableName();
-    m_stringBuilder.append("thread ", m_typeNamer.mangledNameForType(variableReference.resolvedType()), "* ", pointerName, " = &", iterator->value, ";\n");
+    m_stringBuilder.append(m_indent, "thread ", m_typeNamer.mangledNameForType(variableReference.resolvedType()), "* ", pointerName, " = &", iterator->value, ";\n");
     appendLeftValue(variableReference, iterator->value, pointerName, Nullability::NotNull);
 }
 
index 0fa15c8..faef60e 100644 (file)
@@ -120,7 +120,7 @@ static const char* vectorSuffix(int vectorLength)
     }
 }
 
-void inlineNativeFunction(StringBuilder& stringBuilder, AST::NativeFunctionDeclaration& nativeFunctionDeclaration, MangledVariableName returnName, const Vector<MangledVariableName>& args, Intrinsics& intrinsics, TypeNamer& typeNamer)
+void inlineNativeFunction(StringBuilder& stringBuilder, AST::NativeFunctionDeclaration& nativeFunctionDeclaration, MangledVariableName returnName, const Vector<MangledVariableName>& args, Intrinsics& intrinsics, TypeNamer& typeNamer, std::function<MangledVariableName()>&& generateNextVariableName, Indentation<4> indent)
 {
     auto asMatrixType = [&] (AST::UnnamedType& unnamedType) -> AST::NativeTypeDeclaration* {
         auto& realType = unnamedType.unifyNode();
@@ -139,37 +139,43 @@ void inlineNativeFunction(StringBuilder& stringBuilder, AST::NativeFunctionDecla
         auto metalReturnTypeName = typeNamer.mangledNameForType(returnType);
 
         if (!nativeFunctionDeclaration.parameters().size()) {
-            stringBuilder.append(returnName, " = { };\n");
+            stringBuilder.append(indent, returnName, " = { };\n");
             return;
         }
 
-
         if (nativeFunctionDeclaration.parameters().size() == 1) {
             auto& parameterType = *nativeFunctionDeclaration.parameters()[0]->type();
             auto metalParameterTypeName = typeNamer.mangledNameForType(parameterType);
-            stringBuilder.append("{\n", metalParameterTypeName, " x = ", args[0], ";\n");
-
-            {
-                auto isEnumerationDefinition = [] (auto& type) {
-                    return is<AST::NamedType>(type) && is<AST::EnumerationDefinition>(downcast<AST::NamedType>(type));
-                };
-                auto& unifiedReturnType = returnType.unifyNode();
-                if (isEnumerationDefinition(unifiedReturnType) && !isEnumerationDefinition(parameterType.unifyNode())) { 
-                    auto& enumerationDefinition = downcast<AST::EnumerationDefinition>(downcast<AST::NamedType>(unifiedReturnType));
-                    stringBuilder.append("    switch (x) {\n");
+            auto variableName = generateNextVariableName();
+
+            stringBuilder.append(indent, metalParameterTypeName, ' ', variableName, " = ", args[0], ";\n");
+
+            auto isEnumerationDefinition = [] (auto& type) {
+                return is<AST::NamedType>(type) && is<AST::EnumerationDefinition>(downcast<AST::NamedType>(type));
+            };
+            auto& unifiedReturnType = returnType.unifyNode();
+            if (isEnumerationDefinition(unifiedReturnType) && !isEnumerationDefinition(parameterType.unifyNode())) {
+                auto& enumerationDefinition = downcast<AST::EnumerationDefinition>(downcast<AST::NamedType>(unifiedReturnType));
+                stringBuilder.append(indent, "switch (", variableName, ") {\n");
+                {
+                    IndentationScope switchScope(indent);
                     bool hasZeroCase = false;
                     for (auto& member : enumerationDefinition.enumerationMembers()) {
                         hasZeroCase |= !member.get().value();
-                        stringBuilder.append("        case ", member.get().value(), ": break;\n");
+                        stringBuilder.append(
+                            indent, "case ", member.get().value(), ":\n",
+                            indent, "    break;\n");
                     }
                     ASSERT_UNUSED(hasZeroCase, hasZeroCase);
-                    stringBuilder.append("        default: x = 0; break; }\n");
+                    stringBuilder.append(
+                        indent, "default:\n",
+                        indent, "    ", variableName, " = 0;\n",
+                        indent, "    break;\n",
+                        indent, "}\n");
                 }
             }
 
-            stringBuilder.append(
-                returnName, " = static_cast<", metalReturnTypeName, ">(x);\n}\n");
-
+            stringBuilder.append(indent, returnName, " = static_cast<", metalReturnTypeName, ">(", variableName, ");\n");
             return;
         }
 
@@ -178,32 +184,33 @@ void inlineNativeFunction(StringBuilder& stringBuilder, AST::NativeFunctionDecla
             unsigned numColumns = matrixType->numberOfMatrixColumns();
             RELEASE_ASSERT(nativeFunctionDeclaration.parameters().size() == numRows || nativeFunctionDeclaration.parameters().size() == numRows * numColumns);
 
-            stringBuilder.append("{\n", metalReturnTypeName, " x;\n");
+            auto variableName = generateNextVariableName();
+
+            stringBuilder.append(indent, metalReturnTypeName, ' ', variableName, ";\n");
 
             // We need to abide by the memory layout we use for matrices here.
             if (nativeFunctionDeclaration.parameters().size() == numRows) {
                 // operator matrixMxN (vectorN, ..., vectorN)
                 for (unsigned i = 0; i < numRows; ++i) {
                     for (unsigned j = 0; j < numColumns; ++j)
-                        stringBuilder.append("x[", j * numRows + i, "] = ", args[i], "[", j, "];\n");
+                        stringBuilder.append(indent, variableName, "[", j * numRows + i, "] = ", args[i], "[", j, "];\n");
                 }
-
             } else {
                 // operator matrixMxN (scalar, ..., scalar)
                 unsigned index = 0;
                 for (unsigned i = 0; i < numRows; ++i) {
                     for (unsigned j = 0; j < numColumns; ++j) {
-                        stringBuilder.append("x[", j * numRows + i, "] = ", args[index], ";\n");
+                        stringBuilder.append(indent, variableName, '[', j * numRows + i, "] = ", args[index], ";\n");
                         ++index;
                     }
                 }
             }
 
-            stringBuilder.append(returnName, " = x;\n}\n");
+            stringBuilder.append(indent, returnName, " = ", variableName, ";\n");
             return;
         }
 
-        stringBuilder.append(returnName, " = ", metalReturnTypeName, "(");
+        stringBuilder.append(indent, returnName, " = ", metalReturnTypeName, "(");
         for (unsigned i = 0; i < nativeFunctionDeclaration.parameters().size(); ++i) {
             if (i > 0)
                 stringBuilder.append(", ");
@@ -222,13 +229,13 @@ void inlineNativeFunction(StringBuilder& stringBuilder, AST::NativeFunctionDecla
         if (is<AST::ArrayType>(unnamedParameterType)) {
             auto& arrayParameterType = downcast<AST::ArrayType>(unnamedParameterType);
             stringBuilder.append(
-                returnName, " = ", arrayParameterType.numElements(), ";\n");
+                indent, returnName, " = ", arrayParameterType.numElements(), ";\n");
             return;
         }
 
         ASSERT(is<AST::ArrayReferenceType>(unnamedParameterType));
         stringBuilder.append(
-            returnName, " = ", args[0], ".length;\n");
+            indent, returnName, " = ", args[0], ".length;\n");
         return;
     }
 
@@ -253,8 +260,8 @@ void inlineNativeFunction(StringBuilder& stringBuilder, AST::NativeFunctionDecla
             fieldName = fieldName.substring(0, fieldName.length() - 1);
 
             stringBuilder.append(
-                returnName, " = ", args[0], ";\n",
-                returnName, '.');
+                indent, returnName, " = ", args[0], ";\n",
+                indent, returnName, '.');
             appendMangledFieldName(fieldName);
             stringBuilder.append(" = ", args[1], ";\n");
 
@@ -264,7 +271,7 @@ void inlineNativeFunction(StringBuilder& stringBuilder, AST::NativeFunctionDecla
         ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
         auto fieldName = nativeFunctionDeclaration.name().substring("operator."_str.length());
         stringBuilder.append(
-            returnName, " = ", args[0], '.');
+            indent, returnName, " = ", args[0], '.');
         appendMangledFieldName(fieldName);
         stringBuilder.append(";\n");
         return;
@@ -275,7 +282,7 @@ void inlineNativeFunction(StringBuilder& stringBuilder, AST::NativeFunctionDecla
         auto fieldName = nativeFunctionDeclaration.name().substring("operator&."_str.length());
 
         stringBuilder.append(
-            returnName, " = &(", args[0], "->");
+            indent, returnName, " = &(", args[0], "->");
 
         auto& unnamedType = *nativeFunctionDeclaration.parameters()[0]->type();
         auto& unifyNode = downcast<AST::PointerType>(unnamedType).elementType().unifyNode();
@@ -298,7 +305,7 @@ void inlineNativeFunction(StringBuilder& stringBuilder, AST::NativeFunctionDecla
         ASSERT(is<AST::ArrayReferenceType>(*nativeFunctionDeclaration.parameters()[0]->type()));
 
         stringBuilder.append(
-            returnName, " = (", args[1], " < ", args[0], ".length) ? ", " &(", args[0], ".pointer[", args[1], "]) : nullptr;\n");
+            indent, returnName, " = (", args[1], " < ", args[0], ".length) ? ", " &(", args[0], ".pointer[", args[1], "]) : nullptr;\n");
             
         return;
     }
@@ -327,35 +334,47 @@ void inlineNativeFunction(StringBuilder& stringBuilder, AST::NativeFunctionDecla
 
             unsigned numberOfRows = getMatrixType().numberOfMatrixRows();
             unsigned numberOfColumns = getMatrixType().numberOfMatrixColumns();
-            stringBuilder.append("do {\n", metalReturnName, " result;\n");
-
-            stringBuilder.append(
-                "    if (", args[1], " >= ", numberOfRows, ") {", returnName, " = ", metalReturnName, "(0); break;}\n",
-                "    result[0] = ", args[0], '[', args[1], "];\n",
-                "    result[1] = ", args[0], '[', args[1], " + ", numberOfRows, "];\n");
-
-            if (numberOfColumns >= 3)
-                stringBuilder.append("    result[2] = ", args[0], '[', args[1], " + ", numberOfRows * 2, "];\n");
-            if (numberOfColumns >= 4)
-                stringBuilder.append("    result[3] = ", args[0], '[', args[1], " + ", numberOfRows * 3, "];\n");
+            
+            stringBuilder.append(indent, "do {\n");
+            {
+                IndentationScope scope(indent);
 
-            stringBuilder.append(
-                "    ", returnName, " = result;\n",
-                "} while (0);\n");
+                stringBuilder.append(
+                    indent, metalReturnName, " result;\n",
+                    indent, "if (", args[1], " >= ", numberOfRows, ") {\n",
+                    indent, "    ", returnName, " = ", metalReturnName, "(0);\n",
+                    indent, "    break;\n",
+                    indent, "}\n",
+                    indent, "result[0] = ", args[0], '[', args[1], "];\n",
+                    indent, "result[1] = ", args[0], '[', args[1], " + ", numberOfRows, "];\n");
+
+                if (numberOfColumns >= 3)
+                    stringBuilder.append(indent, "result[2] = ", args[0], '[', args[1], " + ", numberOfRows * 2, "];\n");
+                if (numberOfColumns >= 4)
+                    stringBuilder.append(indent, "result[3] = ", args[0], '[', args[1], " + ", numberOfRows * 3, "];\n");
+    
+                stringBuilder.append(indent, returnName, " = result;\n");
+            }
+            stringBuilder.append("} while (0);\n");
         } else {
             RELEASE_ASSERT(numTypeArguments == 2);
             unsigned numElements = vectorSize();
 
             auto metalReturnName = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
 
-            stringBuilder.append("do {\n", metalReturnName, " result;\n");
-
-            stringBuilder.append(
-                "    if (", args[1], " >= ", numElements, ") {", returnName, " = ", metalReturnName, "(0); break;}\n",
-                "    result = ", args[0], "[", args[1], "];\n");
-            stringBuilder.append(
-                "    ", returnName, " = result;\n",
-                "} while (0);\n");
+            stringBuilder.append(indent, "do {\n");
+            {
+                IndentationScope scope(indent);
+                stringBuilder.append(
+                    indent, metalReturnName, " result;\n",
+                    indent, "if (", args[1], " >= ", numElements, ") {\n",
+                    indent, "    ", returnName, " = ", metalReturnName, "(0);\n",
+                    indent, "    break;\n",
+                    indent, "}\n",
+                    indent, "result = ", args[0], "[", args[1], "];\n",
+                    indent, returnName, " = result;\n");
+            }
+            stringBuilder.append(indent, "} while (0);\n");
         }
 
         return;
@@ -372,20 +391,27 @@ void inlineNativeFunction(StringBuilder& stringBuilder, AST::NativeFunctionDecla
             unsigned numberOfRows = getMatrixType().numberOfMatrixRows();
             unsigned numberOfColumns = getMatrixType().numberOfMatrixColumns();
 
-            stringBuilder.append("do {\n", metalReturnName, " m = ", args[0], ";\n",
-                metalParameter2Name, " i = ", args[1], ";\n");
+            stringBuilder.append(indent, "do {\n");
+            {
+                IndentationScope scope(indent);
 
-            stringBuilder.append(
-                "    if (i >= ", numberOfRows, ") {", returnName, " = m;\nbreak;}\n",
-                "    m[i] = ", args[2], "[0];\n",
-                "    m[i + ", numberOfRows, "] = ", args[2], "[1];\n");
-            if (numberOfColumns >= 3)
-                stringBuilder.append("    m[i + ", numberOfRows * 2, "] = ", args[2], "[2];\n");
-            if (numberOfColumns >= 4)
-                stringBuilder.append("    m[i + ", numberOfRows * 3, "] = ", args[2], "[3];\n");
-            stringBuilder.append(
-                "    ", returnName, " = m;\n",
-                "} while(0);\n");
+                stringBuilder.append(
+                    indent, metalReturnName, " m = ", args[0], ";\n",
+                    indent, metalParameter2Name, " i = ", args[1], ";\n",
+                    indent, "if (i >= ", numberOfRows, ") {\n",
+                    indent, "    ", returnName, " = m;\n",
+                    indent, "    break;\n",
+                    indent, "}\n",
+                    indent, "m[i] = ", args[2], "[0];\n",
+                    indent, "m[i + ", numberOfRows, "] = ", args[2], "[1];\n");
+                if (numberOfColumns >= 3)
+                    stringBuilder.append(indent, "m[i + ", numberOfRows * 2, "] = ", args[2], "[2];\n");
+                if (numberOfColumns >= 4)
+                    stringBuilder.append(indent, "m[i + ", numberOfRows * 3, "] = ", args[2], "[3];\n");
+                stringBuilder.append(indent, returnName, " = m;\n");
+            }
+            
+            stringBuilder.append(indent, "} while(0);\n");
         } else {
             RELEASE_ASSERT(numTypeArguments == 2);
 
@@ -395,15 +421,21 @@ void inlineNativeFunction(StringBuilder& stringBuilder, AST::NativeFunctionDecla
 
             unsigned numElements = vectorSize();
 
-            stringBuilder.append("do {\n", metalReturnName, " v = ", args[0], ";\n",
-                metalParameter2Name, " i = ", args[1], ";\n");
+            stringBuilder.append(indent, "do {\n");
+            {
+                IndentationScope scope(indent);
 
-            stringBuilder.append(
-                "    if (i >= ", numElements, ") {", returnName, " = v;\nbreak;}\n",
-                "    v[i] = ", args[2], ";\n");
-            stringBuilder.append(
-                "    ", returnName, " = v;\n",
-                "} while(0);\n");
+                stringBuilder.append(
+                    indent, metalReturnName, " v = ", args[0], ";\n",
+                    indent, metalParameter2Name, " i = ", args[1], ";\n",
+                    indent, "if (i >= ", numElements, ") {\n",
+                    indent, "    ", returnName, " = v;\n",
+                    indent, "    break;\n",
+                    indent, "}\n",
+                    indent, "v[i] = ", args[2], ";\n",
+                    indent, returnName, " = v;\n");
+            }
+            stringBuilder.append(indent, "} while(0);\n");
         }
 
         return;
@@ -414,14 +446,25 @@ void inlineNativeFunction(StringBuilder& stringBuilder, AST::NativeFunctionDecla
         auto metalReturnType = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
         if (nativeFunctionDeclaration.parameters().size() == 1) {
             if (auto* matrixType = asMatrixType(nativeFunctionDeclaration.type())) {
-                stringBuilder.append(
-                    "{\n", metalReturnType, " x = ", args[0], ";\n",
-                    "for (size_t i = 0; i < x.size(); ++i) x[i] = ", operatorName, "x[i];\n",
-                    returnName, " = x;\n}\n");
+                stringBuilder.append(indent, "{\n");
+                {
+                    IndentationScope scope(indent);
+                    stringBuilder.append(
+                        indent, metalReturnType, " x = ", args[0], ";\n",
+                        indent, "for (size_t i = 0; i < x.size(); ++i)\n",
+                        indent, "    x[i] = ", operatorName, "x[i];\n",
+                        indent, returnName, " = x;\n");
+                }
+                stringBuilder.append(indent, "}\n");
             } else {
-                stringBuilder.append(
-                    "{\n", metalReturnType, " x = ", args[0], ";\n", 
-                    returnName, " = ", operatorName, "x;\n}\n");
+                stringBuilder.append(indent, "{\n");
+                {
+                    IndentationScope scope(indent);
+                    stringBuilder.append(
+                        indent, metalReturnType, " x = ", args[0], ";\n",
+                        indent, returnName, " = ", operatorName, "x;\n");
+                }
+                stringBuilder.append(indent, "}\n");
             }
             return;
         }
@@ -430,31 +473,49 @@ void inlineNativeFunction(StringBuilder& stringBuilder, AST::NativeFunctionDecla
         if (auto* leftMatrix = asMatrixType(*nativeFunctionDeclaration.parameters()[0]->type())) {
             if (auto* rightMatrix = asMatrixType(*nativeFunctionDeclaration.parameters()[1]->type())) {
                 // matrix <op> matrix
-                stringBuilder.append(
-                    "{\n", metalReturnType, " x;\n",
-                    "for (size_t i = 0; i < x.size(); ++i) x[i] = ", args[0], "[i] ", operatorName, ' ', args[1], "[i];\n",
-                    returnName, " = x;\n}\n");
+                stringBuilder.append(indent, "{\n");
+                {
+                    IndentationScope scope(indent);
+                    stringBuilder.append(
+                        indent, metalReturnType, " x;\n",
+                        indent, "for (size_t i = 0; i < x.size(); ++i)\n",
+                        indent, "    x[i] = ", args[0], "[i] ", operatorName, ' ', args[1], "[i];\n",
+                        indent, returnName, " = x;\n");
+                }
+                stringBuilder.append(indent, "}\n");
             } else {
                 // matrix <op> scalar
-                stringBuilder.append(
-                    "{\n", metalReturnType, " x;\n",
-                    "for (size_t i = 0; i < x.size(); ++i) x[i] = ", args[0], "[i] ", operatorName, ' ', args[1], ";\n",
-                    returnName, " = x;\n}\n");
+                stringBuilder.append(indent, "{\n");
+                {
+                    IndentationScope scope(indent);
+                    stringBuilder.append(
+                        indent, metalReturnType, " x;\n",
+                        indent, "for (size_t i = 0; i < x.size(); ++i)\n",
+                        indent, "    x[i] = ", args[0], "[i] ", operatorName, ' ', args[1], ";\n",
+                        indent, returnName, " = x;\n");
+                }
+                stringBuilder.append(indent, "}\n");
             }
         } else if (auto* rightMatrix = asMatrixType(*nativeFunctionDeclaration.parameters()[1]->type())) {
             ASSERT(!asMatrixType(*nativeFunctionDeclaration.parameters()[0]->type()));
             // scalar <op> matrix
-            stringBuilder.append(
-                "{\n", metalReturnType, " x;\n",
-                "for (size_t i = 0; i < x.size(); ++i) x[i] = ", args[0], ' ', operatorName, ' ', args[1], "[i];\n",
-                returnName, " = x;\n}\n");
+            stringBuilder.append(indent, "{\n");
+            {
+                IndentationScope scope(indent);
+                stringBuilder.append(
+                    indent, metalReturnType, " x;\n",
+                    indent, "for (size_t i = 0; i < x.size(); ++i)\n",
+                    indent, "    x[i] = ", args[0], ' ', operatorName, ' ', args[1], "[i];\n",
+                    indent, returnName, " = x;\n");
+            }
+            stringBuilder.append(indent, "}\n");
         } else {
             // scalar <op> scalar
             // vector <op> vector
             // vector <op> scalar
             // scalar <op> vector
             stringBuilder.append(
-                returnName, " = ", args[0], ' ', operatorName, ' ', args[1], ";\n");
+                indent, returnName, " = ", args[0], ' ', operatorName, ' ', args[1], ";\n");
         }
 
         return;
@@ -487,14 +548,14 @@ void inlineNativeFunction(StringBuilder& stringBuilder, AST::NativeFunctionDecla
         || nativeFunctionDeclaration.name() == "length") {
         ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
         stringBuilder.append(
-            returnName, " = ", mapFunctionName(nativeFunctionDeclaration.name()), '(', args[0], ");\n");
+            indent, returnName, " = ", mapFunctionName(nativeFunctionDeclaration.name()), '(', args[0], ");\n");
         return;
     }
 
     if (nativeFunctionDeclaration.name() == "pow" || nativeFunctionDeclaration.name() == "atan2") {
         ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
         stringBuilder.append(
-            returnName, " = ", nativeFunctionDeclaration.name(), "(", args[0], ", ", args[1], ");\n");
+            indent, returnName, " = ", nativeFunctionDeclaration.name(), "(", args[0], ", ", args[1], ");\n");
         return;
     }
 
@@ -502,13 +563,20 @@ void inlineNativeFunction(StringBuilder& stringBuilder, AST::NativeFunctionDecla
         ASSERT(nativeFunctionDeclaration.parameters().size() == 3);
         if (asMatrixType(nativeFunctionDeclaration.type())) {
             auto metalReturnType = typeNamer.mangledNameForType(nativeFunctionDeclaration.type());
-            stringBuilder.append(
-                "{\n", metalReturnType, " x;\n",
-                "for (size_t i = 0; i < x.size(); ++i) x[i] = clamp(", args[0], "[i], ", args[1], "[i], ", args[2], "[i]);",
-                returnName, " = x;\n}\n");
+            
+            stringBuilder.append(indent, "{\n");
+            {
+                IndentationScope scope(indent);
+                stringBuilder.append(
+                    indent, metalReturnType, " x;\n",
+                    indent, "for (size_t i = 0; i < x.size(); ++i) \n",
+                    indent, "    x[i] = clamp(", args[0], "[i], ", args[1], "[i], ", args[2], "[i]);",
+                    indent, returnName, " = x;\n");
+            }
+            stringBuilder.append(indent, "}\n");
         } else {
             stringBuilder.append(
-                returnName, " = clamp(", args[0], ", ", args[1], ", ", args[2], ");\n");
+                indent, returnName, " = clamp(", args[0], ", ", args[1], ", ", args[2], ");\n");
         }
         return;
     }
@@ -516,23 +584,23 @@ void inlineNativeFunction(StringBuilder& stringBuilder, AST::NativeFunctionDecla
     if (nativeFunctionDeclaration.name() == "AllMemoryBarrierWithGroupSync") {
         ASSERT(!nativeFunctionDeclaration.parameters().size());
         stringBuilder.append(
-            "threadgroup_barrier(mem_flags::mem_device);\n"
-            "threadgroup_barrier(mem_flags::mem_threadgroup);\n"
-            "threadgroup_barrier(mem_flags::mem_texture);\n");
+            indent, "threadgroup_barrier(mem_flags::mem_device);\n",
+            indent, "threadgroup_barrier(mem_flags::mem_threadgroup);\n",
+            indent, "threadgroup_barrier(mem_flags::mem_texture);\n");
         return;
     }
 
     if (nativeFunctionDeclaration.name() == "DeviceMemoryBarrierWithGroupSync") {
         ASSERT(!nativeFunctionDeclaration.parameters().size());
         stringBuilder.append(
-            "threadgroup_barrier(mem_flags::mem_device);\n");
+            indent, "threadgroup_barrier(mem_flags::mem_device);\n");
         return;
     }
 
     if (nativeFunctionDeclaration.name() == "GroupMemoryBarrierWithGroupSync") {
         ASSERT(!nativeFunctionDeclaration.parameters().size());
         stringBuilder.append(
-            "threadgroup_barrier(mem_flags::mem_threadgroup);\n");
+            indent, "threadgroup_barrier(mem_flags::mem_threadgroup);\n");
         return;
     }
 
@@ -540,15 +608,15 @@ void inlineNativeFunction(StringBuilder& stringBuilder, AST::NativeFunctionDecla
         if (nativeFunctionDeclaration.name() == "InterlockedCompareExchange") {
             ASSERT(nativeFunctionDeclaration.parameters().size() == 4);
             stringBuilder.append(
-                "atomic_compare_exchange_weak_explicit(", args[0], ", &", args[1], ", ", args[2], ", memory_order_relaxed, memory_order_relaxed);\n",
-                '*', args[3], " = ", args[1], ";\n");
+                indent, "atomic_compare_exchange_weak_explicit(", args[0], ", &", args[1], ", ", args[2], ", memory_order_relaxed, memory_order_relaxed);\n",
+                indent, '*', args[3], " = ", args[1], ";\n");
             return;
         }
 
         ASSERT(nativeFunctionDeclaration.parameters().size() == 3);
         auto name = atomicName(nativeFunctionDeclaration.name().substring("Interlocked"_str.length()));
         stringBuilder.append(
-            '*', args[2], " = atomic_", name, "_explicit(", args[0], ", ", args[1], ", memory_order_relaxed);\n");
+            indent, '*', args[2], " = atomic_", name, "_explicit(", args[0], ", ", args[1], ", memory_order_relaxed);\n");
         return;
     }
 
@@ -562,7 +630,7 @@ void inlineNativeFunction(StringBuilder& stringBuilder, AST::NativeFunctionDecla
         auto returnVectorLength = vectorLength(returnType);
 
         stringBuilder.append(
-            returnName, " = ", args[0], ".sample(", args[1], ", ");
+            indent, returnName, " = ", args[0], ".sample(", args[1], ", ");
 
         if (textureType.isTextureArray()) {
             ASSERT(locationVectorLength > 1);
@@ -589,41 +657,66 @@ void inlineNativeFunction(StringBuilder& stringBuilder, AST::NativeFunctionDecla
         auto returnVectorLength = vectorLength(returnType);
 
         auto metalReturnName = typeNamer.mangledNameForType(returnType);
-        stringBuilder.append("do {\n");
 
-        if (textureType.isTextureArray()) {
-            ASSERT(locationVectorLength > 1);
-            String dimensions[] = { "width"_str, "height"_str, "depth"_str };
-            for (int i = 0; i < locationVectorLength - 1; ++i) {
-                auto suffix = "xyzw"_str.substring(i, 1);
-                stringBuilder.append("    if (", args[1], '.', suffix, " < 0 || static_cast<uint32_t>(", args[1], '.', suffix, ") >= ", args[0], ".get_", dimensions[i], "()) {", returnName, " = ", metalReturnName, "(0); break;}\n");
-            }
-            auto suffix = "xyzw"_str.substring(locationVectorLength - 1, 1);
-            stringBuilder.append("    if (", args[1], '.', suffix, " < 0 || static_cast<uint32_t>(", args[1], '.', suffix, ") >= ", args[0], ".get_array_size()) {", returnName, " = ", metalReturnName, "(0); break;}\n");
-        } else {
-            if (locationVectorLength == 1)
-                stringBuilder.append("    if (", args[1], " < 0 || static_cast<uint32_t>(", args[1], ") >= ", args[0], ".get_width()) { ", returnName, " = ", metalReturnName, "(0); break;}\n");
-            else {
+        stringBuilder.append(indent, "do {\n");
+        {
+            IndentationScope scope(indent);
+
+            if (textureType.isTextureArray()) {
+                ASSERT(locationVectorLength > 1);
+                String dimensions[] = { "width"_str, "height"_str, "depth"_str };
+                for (int i = 0; i < locationVectorLength - 1; ++i) {
+                    auto suffix = "xyzw"_str.substring(i, 1);
+                    stringBuilder.append(
+                        indent, "if (", args[1], '.', suffix, " < 0 || static_cast<uint32_t>(", args[1], '.', suffix, ") >= ", args[0], ".get_", dimensions[i], "()) {\n",
+                        indent, "    ", returnName, " = ", metalReturnName, "(0);\n",
+                        indent, "    break;\n",
+                        indent, "}\n");
+                }
+                auto suffix = "xyzw"_str.substring(locationVectorLength - 1, 1);
                 stringBuilder.append(
-                    "    if (", args[1], ".x < 0 || static_cast<uint32_t>(", args[1], ".x) >= ", args[0], ".get_width()) {", returnName, " = ", metalReturnName, "(0); break;}\n"
-                    "    if (", args[1], ".y < 0 || static_cast<uint32_t>(", args[1], ".y) >= ", args[0], ".get_height()) {", returnName, " = ", metalReturnName, "(0); break;}\n");
-
-                if (locationVectorLength >= 3)
-                    stringBuilder.append("    if (", args[1], ".z < 0 || static_cast<uint32_t>(", args[1], ".z) >= ", args[0], ".get_depth()) {", returnName, " = ", metalReturnName, "(0); break;}\n");
+                    indent, "if (", args[1], '.', suffix, " < 0 || static_cast<uint32_t>(", args[1], '.', suffix, ") >= ", args[0], ".get_array_size()) {\n",
+                    indent, "    ", returnName, " = ", metalReturnName, "(0);\n",
+                    indent, "    break;\n",
+                    indent, "}\n");
+            } else {
+                if (locationVectorLength == 1) {
+                    stringBuilder.append(
+                        indent, "if (", args[1], " < 0 || static_cast<uint32_t>(", args[1], ") >= ", args[0], ".get_width()) {\n",
+                        indent, "    ", returnName, " = ", metalReturnName, "(0);\n",
+                        indent, "    break;\n",
+                        indent, "}\n");
+                } else {
+                    stringBuilder.append(
+                        indent, "if (", args[1], ".x < 0 || static_cast<uint32_t>(", args[1], ".x) >= ", args[0], ".get_width()) {\n",
+                        indent, "    ", returnName, " = ", metalReturnName, "(0);\n",
+                        indent, "    break;\n",
+                        indent, "}\n",
+                        indent, "if (", args[1], ".y < 0 || static_cast<uint32_t>(", args[1], ".y) >= ", args[0], ".get_height()) {\n",
+                        indent, "    ", returnName, " = ", metalReturnName, "(0);\n",
+                        indent, "    break;\n",
+                        indent, "}\n");
+                    if (locationVectorLength >= 3) {
+                        stringBuilder.append(
+                            indent, "if (", args[1], ".z < 0 || static_cast<uint32_t>(", args[1], ".z) >= ", args[0], ".get_depth()) {\n",
+                            indent, "    ", returnName, " = ", metalReturnName, "(0);\n",
+                            indent, "    break;\n",
+                            indent, "}\n");
+                    }
+                }
             }
+            stringBuilder.append(indent, returnName, " = ", args[0], ".read(");
+            if (textureType.isTextureArray()) {
+                ASSERT(locationVectorLength > 1);
+                stringBuilder.append("uint", vectorSuffix(locationVectorLength - 1), '(', args[1], '.', "xyzw"_str.substring(0, locationVectorLength - 1), "), uint(", args[1], '.', "xyzw"_str.substring(locationVectorLength - 1, 1), ')');
+            } else
+                stringBuilder.append("uint", vectorSuffix(locationVectorLength), '(', args[1], ')');
+            stringBuilder.append(')');
+            if (!textureType.isDepthTexture())
+                stringBuilder.append('.', "xyzw"_str.substring(0, returnVectorLength));
+            stringBuilder.append(";\n");
         }
-        stringBuilder.append("    ", returnName, " = ", args[0], ".read(");
-        if (textureType.isTextureArray()) {
-            ASSERT(locationVectorLength > 1);
-            stringBuilder.append("uint", vectorSuffix(locationVectorLength - 1), '(', args[1], '.', "xyzw"_str.substring(0, locationVectorLength - 1), "), uint(", args[1], '.', "xyzw"_str.substring(locationVectorLength - 1, 1), ')');
-        } else
-            stringBuilder.append("uint", vectorSuffix(locationVectorLength), '(', args[1], ')');
-        stringBuilder.append(')');
-        if (!textureType.isDepthTexture())
-            stringBuilder.append('.', "xyzw"_str.substring(0, returnVectorLength));
-        stringBuilder.append(
-            ";\n"
-            "} while(0);\n");
+        stringBuilder.append(indent, "} while(0);\n");
 
         return;
     }
@@ -631,14 +724,14 @@ void inlineNativeFunction(StringBuilder& stringBuilder, AST::NativeFunctionDecla
     if (nativeFunctionDeclaration.name() == "load") {
         ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
         stringBuilder.append(
-            returnName, " = atomic_load_explicit(", args[0], ", memory_order_relaxed);\n");
+            indent, returnName, " = atomic_load_explicit(", args[0], ", memory_order_relaxed);\n");
         return;
     }
 
     if (nativeFunctionDeclaration.name() == "store") {
         ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
         stringBuilder.append(
-            "atomic_store_explicit(", args[0], ", ", args[1], ", memory_order_relaxed);\n");
+            indent, "atomic_store_explicit(", args[0], ", ", args[1], ", memory_order_relaxed);\n");
         return;
     }
 
@@ -674,37 +767,37 @@ void inlineNativeFunction(StringBuilder& stringBuilder, AST::NativeFunctionDecla
         ASSERT(index == nativeFunctionDeclaration.parameters().size());
 
         stringBuilder.append(
-            "if (", widthName, ")\n"
-            "    *", widthName, " = ", args[0], ".get_width(");
-
+            indent, "if (", widthName, ")\n",
+            indent, "    *", widthName, " = ", args[0], ".get_width(");
         if (hasMipLevel)
             stringBuilder.append(args[1]);
         stringBuilder.append(");\n");
+
         if (heightName) {
             stringBuilder.append(
-                "    if (", *heightName, ")\n"
-                "        *", *heightName, " = ", args[0], ".get_height(");
+                indent, "if (", *heightName, ")\n",
+                indent, "    *", *heightName, " = ", args[0], ".get_height(");
             if (hasMipLevel)
                 stringBuilder.append(args[1]);
             stringBuilder.append(");\n");
         }
         if (depthName) {
             stringBuilder.append(
-                "    if (", *depthName, ")\n"
-                "        *", *depthName, " = ", args[0], ".get_depth(");
+                indent, "if (", *depthName, ")\n",
+                indent, "    *", *depthName, " = ", args[0], ".get_depth(");
             if (hasMipLevel)
                 stringBuilder.append(args[1]);
             stringBuilder.append(");\n");
         }
         if (elementsName) {
             stringBuilder.append(
-                "    if (", *elementsName, ")\n"
-                "        *", *elementsName, " = ", args[0], ".get_array_size();\n");
+                indent, "if (", *elementsName, ")\n",
+                indent, "    *", *elementsName, " = ", args[0], ".get_array_size();\n");
         }
         if (numberOfLevelsName) {
             stringBuilder.append(
-                "    if (", *numberOfLevelsName, ")\n"
-                "        *", *numberOfLevelsName, " = ", args[0], ".get_num_mip_levels();\n");
+                indent, "if (", *numberOfLevelsName, ")\n",
+                indent, "    *", *numberOfLevelsName, " = ", args[0], ".get_num_mip_levels();\n");
         }
         return;
     }
@@ -757,38 +850,52 @@ void inlineNativeFunction(StringBuilder& stringBuilder, AST::NativeFunctionDecla
         auto metalInnerTypeName = typeNamer.mangledNameForType(itemVectorInnerType);
 
         stringBuilder.append("do {\n");
-        if (textureType.isTextureArray()) {
-            ASSERT(locationVectorLength > 1);
-            String dimensions[] = { "width"_str, "height"_str, "depth"_str };
-            for (int i = 0; i < locationVectorLength - 1; ++i) {
-                auto suffix = "xyzw"_str.substring(i, 1);
-                stringBuilder.append("    if (", args[2], ".", suffix, " >= ", args[0], ".get_", dimensions[i], "()) break;\n");
-            }
-            auto suffix = "xyzw"_str.substring(locationVectorLength - 1, 1);
-            stringBuilder.append("    if (", args[2], '.', suffix, " >= ", args[0], ".get_array_size()) break;\n");
-        } else {
-            if (locationVectorLength == 1)
-                stringBuilder.append("    if (", args[2], " >= ", args[0], ".get_width()) break;\n");
-            else {
+        {
+            IndentationScope scope(indent);
+
+            if (textureType.isTextureArray()) {
+                ASSERT(locationVectorLength > 1);
+                String dimensions[] = { "width"_str, "height"_str, "depth"_str };
+                for (int i = 0; i < locationVectorLength - 1; ++i) {
+                    auto suffix = "xyzw"_str.substring(i, 1);
+                    stringBuilder.append(
+                        indent, "if (", args[2], ".", suffix, " >= ", args[0], ".get_", dimensions[i], "())\n",
+                        indent, "    break;\n");
+                }
+                auto suffix = "xyzw"_str.substring(locationVectorLength - 1, 1);
                 stringBuilder.append(
-                    "    if (", args[2], ".x >= ", args[0], ".get_width()) break;\n"
-                    "    if (", args[2], ".y >= ", args[0], ".get_height()) break;\n");
-                if (locationVectorLength >= 3)
-                    stringBuilder.append("    if (", args[2], ".z >= ", args[0], ".get_depth()) break;\n");
+                    indent, "if (", args[2], '.', suffix, " >= ", args[0], ".get_array_size())\n",
+                    indent, "    break;\n");
+            } else {
+                if (locationVectorLength == 1) {
+                    stringBuilder.append(
+                        indent, "if (", args[2], " >= ", args[0], ".get_width()) \n",
+                        indent, "    break;\n");
+                } else {
+                    stringBuilder.append(
+                        indent, "if (", args[2], ".x >= ", args[0], ".get_width())\n",
+                        indent, "    break;\n",
+                        indent, "if (", args[2], ".y >= ", args[0], ".get_height())\n",
+                        indent, "    break;\n");
+                    if (locationVectorLength >= 3) {
+                        stringBuilder.append(
+                            indent, "if (", args[2], ".z >= ", args[0], ".get_depth())\n",
+                            indent, "    break;\n");
+                    }
+                }
             }
+            stringBuilder.append(indent, args[0], ".write(vec<", metalInnerTypeName, ", 4>(", args[1]);
+            for (int i = 0; i < 4 - itemVectorLength; ++i)
+                stringBuilder.append(", 0");
+            stringBuilder.append("), ");
+            if (textureType.isTextureArray()) {
+                ASSERT(locationVectorLength > 1);
+                stringBuilder.append("uint", vectorSuffix(locationVectorLength - 1), '(', args[2], '.', "xyzw"_str.substring(0, locationVectorLength - 1), "), uint(", args[2], ".", "xyzw"_str.substring(locationVectorLength - 1, 1), ')');
+            } else
+                stringBuilder.append("uint", vectorSuffix(locationVectorLength), '(', args[2], ')');
+            stringBuilder.append(");\n");
         }
-        stringBuilder.append("    ", args[0], ".write(vec<", metalInnerTypeName, ", 4>(", args[1]);
-        for (int i = 0; i < 4 - itemVectorLength; ++i)
-            stringBuilder.append(", 0");
-        stringBuilder.append("), ");
-        if (textureType.isTextureArray()) {
-            ASSERT(locationVectorLength > 1);
-            stringBuilder.append("uint", vectorSuffix(locationVectorLength - 1), '(', args[2], '.', "xyzw"_str.substring(0, locationVectorLength - 1), "), uint(", args[2], ".", "xyzw"_str.substring(locationVectorLength - 1, 1), ')');
-        } else
-            stringBuilder.append("uint", vectorSuffix(locationVectorLength), '(', args[2], ')');
-        stringBuilder.append(
-            ");\n"
-            "} while(0);\n");
+        stringBuilder.append(indent, "} while(0);\n");
 
         return;
     }
index 18e6f5e..483e90a 100644 (file)
@@ -45,7 +45,7 @@ namespace Metal {
 
 class TypeNamer;
 
-void inlineNativeFunction(StringBuilder&, AST::NativeFunctionDeclaration&, MangledVariableName returnName, const Vector<MangledVariableName>& argumentNames, Intrinsics&, TypeNamer&);
+void inlineNativeFunction(StringBuilder&, AST::NativeFunctionDeclaration&, MangledVariableName returnName, const Vector<MangledVariableName>& argumentNames, Intrinsics&, TypeNamer&, std::function<MangledVariableName()>&&, Indentation<4>);
 
 }
 
index 08f951e..ce3144b 100644 (file)
@@ -375,7 +375,7 @@ void TypeNamer::emitUnnamedTypeDefinition(StringBuilder& stringBuilder, BaseType
         auto& arrayReferenceType = downcast<ArrayReferenceTypeNameNode>(baseTypeNameNode);
         ASSERT(baseTypeNameNode.parent());
         stringBuilder.append(
-            "struct ", arrayReferenceType.mangledName(), "\n"
+            "struct ", arrayReferenceType.mangledName(), " {\n"
             "    ", toString(arrayReferenceType.addressSpace()), ' ', arrayReferenceType.parent()->mangledName(), "* pointer;\n"
             "    uint32_t length;\n"
             "};\n"