Attachment 359646 Details for Bug 193531 – Patch

[patch] Patch

patch.txt (text/plain), 120.43 KB, created by Myles C. Maxfield on 2019-01-20 13:37:05 PST

(hide)

Description:

Filename:

MIME Type:

Creator: Myles C. Maxfield

Created: 2019-01-20 13:37:05 PST

Size: 120.43 KB

patch

obsolete

>commit 9f81dd6ff79355cf5fef8a3120d9de006d746971
>Author: Myles C. Maxfield <mmaxfield@apple.com>
>Date:   Fri Jan 18 17:49:52 2019 -0800
>
>    Metal
>
>diff --git a/Source/WebCore/ChangeLog b/Source/WebCore/ChangeLog
>index 03d66934bc0..9fc904b7519 100644
>--- a/Source/WebCore/ChangeLog
>+++ b/Source/WebCore/ChangeLog
>@@ -1,20 +1,137 @@
> 2019-01-20  Myles C. Maxfield  <mmaxfield@apple.com>
> 
>-        [WHLSL] Add the statement behavior checker
>-        https://bugs.webkit.org/show_bug.cgi?id=193487
>+        [WHLSL] Implement Metal code generation
>+        https://bugs.webkit.org/show_bug.cgi?id=193531
> 
>         Reviewed by NOBODY (OOPS!).
> 
>-        This is a translation of https://github.com/gpuweb/WHLSL/blob/master/Spec/source/index.rst#typing-statements
>-        into C++. It is meant to replace the ReturnChecker and UnreachableCodeChecker in the reference implementation.
>-
>-        No new tests because it isn't hooked up yet. Not enough of the compiler exists to have any meaningful sort
>-        of test. When enough of the compiler is present, I'll port the reference implementation's test suite.
>+        This implements the majority of the metal code generation piece. There are still a few pieces missing,
>+        that I'll add in follow up patches. There's still enough complexity here that this is worth reviewing
>+        on it's own, though.
>+
>+        This patch includes a few pieces:
>+        - Metal typedefs for every WHLSL type. This analysis is actually pretty interesting, because complex
>+              types depend on their inner types, and the inner types need to be emitted first. Therefore,
>+              this patch implements a topological sort when emitting types. Also, WHLSL types need to be de-
>+              duped because array references are implemented in MSL as structs, and if you have two structs
>+              in MSL with the same contents, those two structs are not equal and cannot be assigned to each
>+              other. So, this patch creates a trie to de-dup all the UnnamedTypes, and implements a
>+              dependency graph which includes both nodes in the trie as well as NamedTypes which don't appear
>+              in the trie.
>+        - WHLSL enumeration code generation
>+        - A name mangler, which ensures that no text from the source program is contained within the result
>+              program
>+        - Full support for expressions. An expression like "y = *x + 7;" would be converted to something like
>+              Type1 variable1 = *x;
>+              Type2 variable2 = 7;
>+              Type3 variable3 = variable1 + variable2;
>+              y = variable3;
>+        - Mostly complete support for control flow. This is tricky because of how we transform WHLSL
>+              expressions into C++ statements. Therefore, things like "for ( ; *x + 7; )" is difficult to
>+              compile, because we can't put the "*x + 7" generated statements into the for loop itself.
>+              Instead, we have to emit this code inside the loop, in all the places that would implicitly run
>+              it. This patch doesn't fully handle this, see below. (If MSL supported lambdas, we could put
>+              the statements into a lambda and do something like "for ( ; theLambda(); )" but MSL doesn't
>+              support lambdas.)
>+
>+        Missing pieces:
>+        - Entry point pack / unpack code
>+        - Support for "continue" (See above regarding control flow)
>+        - Knowing whether or not a switch case should end with break or fallthrough
>+        - Trapping
>+        - Zero filling variables
>+        - Code generation for compiler-generated native functions (this patch supports native functions in the
>+              standard library, texture functions, and HLSL's half <-> int functions.
>+
>+        No new tests because it isn't hooked up yet. As soon as we can do entry point packing and unpacking,
>+        I'll start porting the test suite.
> 
>+        * Modules/webgpu/WHLSL/AST/WHLSLBuiltInSemantic.h:
>+        (WebCore::WHLSL::AST::BuiltInSemantic::targetIndex):
>+        * Modules/webgpu/WHLSL/AST/WHLSLExpression.h:
>+        (WebCore::WHLSL::AST::Expression::resolvedType):
>+        (WebCore::WHLSL::AST::Expression::type): Deleted.
>+        * Modules/webgpu/WHLSL/AST/WHLSLFunctionDeclaration.h:
>+        (WebCore::WHLSL::AST::FunctionDeclaration::name):
>+        * Modules/webgpu/WHLSL/AST/WHLSLStructureDefinition.h:
>+        (WebCore::WHLSL::AST::StructureDefinition::find):
>+        * Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp: Added.
>+        (WebCore::WHLSL::Metal::EntryPointScaffolding::EntryPointScaffolding):
>+        (WebCore::WHLSL::Metal::EntryPointScaffolding::helperTypes):
>+        (WebCore::WHLSL::Metal::EntryPointScaffolding::signature):
>+        (WebCore::WHLSL::Metal::EntryPointScaffolding::unpack):
>+        (WebCore::WHLSL::Metal::EntryPointScaffolding::pack):
>+        * Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.h: Copied from Source/WebCore/Modules/webgpu/WHLSL/WHLSLLoopChecker.h.
>+        * Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.cpp: Added.
>+        (WebCore::WHLSL::Metal::FunctionDeclarationWriter::FunctionDeclarationWriter):
>+        (WebCore::WHLSL::Metal::FunctionDeclarationWriter::toString):
>+        (WebCore::WHLSL::Metal::FunctionDefinitionWriter::FunctionDefinitionWriter):
>+        (WebCore::WHLSL::Metal::FunctionDefinitionWriter::toString):
>+        (WebCore::WHLSL::Metal::FunctionDefinitionWriter::constantExpressionString):
>+        (WebCore::WHLSL::Metal::FunctionDefinitionWriter::generateNextVariableName):
>+        (WebCore::WHLSL::Metal::metalFunctions):
>+        * Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.h: Copied from Source/WebCore/Modules/webgpu/WHLSL/WHLSLLoopChecker.h.
>+        * Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.cpp: Copied from Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLStructureDefinition.h.
>+        (WebCore::WHLSL::Metal::generateMetalCode):
>+        * Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.h: Copied from Source/WebCore/Modules/webgpu/WHLSL/WHLSLLoopChecker.h.
>+        * Modules/webgpu/WHLSL/Metal/WHLSLNativeFunctionWriter.cpp: Added.
>+        (WebCore::WHLSL::Metal::getNativeName):
>+        (WebCore::WHLSL::Metal::mapFunctionName):
>+        (WebCore::WHLSL::Metal::convertAddressSpace):
>+        (WebCore::WHLSL::Metal::writeNativeFunction):
>+        * Modules/webgpu/WHLSL/Metal/WHLSLNativeFunctionWriter.h: Copied from Source/WebCore/Modules/webgpu/WHLSL/WHLSLLoopChecker.h.
>+        * Modules/webgpu/WHLSL/Metal/WHLSLTypeNamer.cpp: Added.
>+        (WebCore::WHLSL::Metal::BaseTypeNameNode::BaseTypeNameNode):
>+        (WebCore::WHLSL::Metal::BaseTypeNameNode::isArrayTypeNameNode const):
>+        (WebCore::WHLSL::Metal::BaseTypeNameNode::isArrayReferenceTypeNameNode const):
>+        (WebCore::WHLSL::Metal::BaseTypeNameNode::isPointerTypeNameNode const):
>+        (WebCore::WHLSL::Metal::BaseTypeNameNode::isReferenceTypeNameNode const):
>+        (WebCore::WHLSL::Metal::BaseTypeNameNode::children):
>+        (WebCore::WHLSL::Metal::BaseTypeNameNode::append):
>+        (WebCore::WHLSL::Metal::BaseTypeNameNode::parent):
>+        (WebCore::WHLSL::Metal::BaseTypeNameNode::mangledName const):
>+        (WebCore::WHLSL::Metal::ArrayTypeNameNode::ArrayTypeNameNode):
>+        (WebCore::WHLSL::Metal::ArrayTypeNameNode::numElements const):
>+        (WebCore::WHLSL::Metal::ArrayReferenceTypeNameNode::ArrayReferenceTypeNameNode):
>+        (WebCore::WHLSL::Metal::ArrayReferenceTypeNameNode::addressSpace const):
>+        (WebCore::WHLSL::Metal::PointerTypeNameNode::PointerTypeNameNode):
>+        (WebCore::WHLSL::Metal::PointerTypeNameNode::addressSpace const):
>+        (WebCore::WHLSL::Metal::ReferenceTypeNameNode::ReferenceTypeNameNode):
>+        (WebCore::WHLSL::Metal::ReferenceTypeNameNode::namedType):
>+        (WebCore::WHLSL::Metal::TypeNamer::TypeNamer):
>+        (WebCore::WHLSL::Metal::TypeNamer::visit):
>+        (WebCore::WHLSL::Metal::findInVector):
>+        (WebCore::WHLSL::Metal::find):
>+        (WebCore::WHLSL::Metal::TypeNamer::mangledNameForType):
>+        (WebCore::WHLSL::Metal::TypeNamer::createNameNode):
>+        (WebCore::WHLSL::Metal::TypeNamer::insert):
>+        (WebCore::WHLSL::Metal::MetalTypeDeclarationWriter::MetalTypeDeclarationWriter):
>+        (WebCore::WHLSL::Metal::MetalTypeDeclarationWriter::toString):
>+        (WebCore::WHLSL::Metal::MetalTypeDeclarationWriter::visit):
>+        (WebCore::WHLSL::Metal::TypeNamer::metalTypeDeclarations):
>+        (WebCore::WHLSL::Metal::toString):
>+        (WebCore::WHLSL::Metal::TypeNamer::emitUnnamedTypeDefinition):
>+        (WebCore::WHLSL::Metal::TypeNamer::emitNamedTypeDefinition):
>+        (WebCore::WHLSL::Metal::TypeNamer::metalTypeDefinitions):
>+        (WebCore::WHLSL::Metal::TypeNamer::mangledNameForEnumerationMember):
>+        (WebCore::WHLSL::Metal::TypeNamer::mangledNameForStructureElement):
>+        (WebCore::WHLSL::Metal::TypeNamer::metalTypes):
>+        * Modules/webgpu/WHLSL/Metal/WHLSLTypeNamer.h: Added.
>+        (WebCore::WHLSL::Metal::TypeNamer::generateNextTypeName):
>+        (WebCore::WHLSL::Metal::TypeNamer::generateNextStructureElementName):
>+        (WebCore::WHLSL::Metal::TypeNamer::generateNextEnumerationMemberName):
>+        * Modules/webgpu/WHLSL/WHLSLChecker.cpp:
>+        (WebCore::WHLSL::checkSemantics):
>+        (WebCore::WHLSL::Checker::visit):
>+        * Modules/webgpu/WHLSL/WHLSLGatherEntryPointItems.cpp:
>+        (WebCore::WHLSL::Gatherer::visit):
>+        * Modules/webgpu/WHLSL/WHLSLGatherEntryPointItems.h:
>+        (WebCore::WHLSL::EntryPointItem::EntryPointItem):
>+        * Modules/webgpu/WHLSL/WHLSLLoopChecker.cpp: Removed.
>         * Modules/webgpu/WHLSL/WHLSLStatementBehaviorChecker.cpp: Added.
>         (WebCore::WHLSL::StatementBehaviorChecker::takeFunctionBehavior):
>         (WebCore::WHLSL::checkStatementBehavior):
>-        * Modules/webgpu/WHLSL/WHLSLStatementBehaviorChecker.h: Added.
>+        * Modules/webgpu/WHLSL/WHLSLStatementBehaviorChecker.h: Renamed from Source/WebCore/Modules/webgpu/WHLSL/WHLSLLoopChecker.h.
>         * Sources.txt:
>         * WebCore.xcodeproj/project.pbxproj:
> 
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLBuiltInSemantic.h b/Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLBuiltInSemantic.h
>index 43b6986af7d..7b36caabc6b 100644
>--- a/Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLBuiltInSemantic.h
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLBuiltInSemantic.h
>@@ -69,6 +69,7 @@ public:
>     BuiltInSemantic(BuiltInSemantic&&) = default;
> 
>     Variable variable() const { return m_variable; }
>+    Optional<unsigned>& targetIndex() { return m_targetIndex; }
> 
>     bool operator==(const BuiltInSemantic& other) const
>     {
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLExpression.h b/Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLExpression.h
>index 8fef62959d7..01f2ffb806a 100644
>--- a/Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLExpression.h
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLExpression.h
>@@ -57,7 +57,7 @@ public:
> 
>     const Lexer::Token& origin() const { return m_origin; }
> 
>-    UnnamedType* type() { return m_type ? &*m_type : nullptr; }
>+    UnnamedType* resolvedType() { return m_type ? &*m_type : nullptr; }
> 
>     void setType(UniqueRef<UnnamedType>&& type)
>     {
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLFunctionDeclaration.h b/Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLFunctionDeclaration.h
>index 50828584607..c15466a70bb 100644
>--- a/Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLFunctionDeclaration.h
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLFunctionDeclaration.h
>@@ -70,6 +70,7 @@ public:
>     const UnnamedType& type() const { return m_type; }
>     UnnamedType& type() { return m_type; }
>     const String& name() const { return m_name; }
>+    String& name() { return m_name; }
>     bool isCast() const { return m_name == "operator cast"; }
>     const VariableDeclarations& parameters() const { return m_parameters; }
>     VariableDeclarations& parameters() { return m_parameters; }
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLStructureDefinition.h b/Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLStructureDefinition.h
>index 456c6e1238b..0260d4506a5 100644
>--- a/Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLStructureDefinition.h
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLStructureDefinition.h
>@@ -55,6 +55,15 @@ public:
>     bool isStructureDefinition() const override { return true; }
> 
>     StructureElements& structureElements() { return m_structureElements; }
>+    StructureElement* find(String& name)
>+    {
>+        auto iterator = std::find_if(m_structureElements.begin(), m_structureElements.end(), [&](StructureElement& structureElement) -> bool {
>+            return structureElement.name() == name;
>+        });
>+        if (iterator == m_structureElements.end())
>+            return nullptr;
>+        return &*iterator;
>+    }
> 
> private:
>     StructureElements m_structureElements;
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp
>new file mode 100644
>index 00000000000..3d1b076392f
>--- /dev/null
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp
>@@ -0,0 +1,77 @@
>+/*
>+ * Copyright (C) 2019 Apple Inc. All rights reserved.
>+ *
>+ * Redistribution and use in source and binary forms, with or without
>+ * modification, are permitted provided that the following conditions
>+ * are met:
>+ * 1. Redistributions of source code must retain the above copyright
>+ *    notice, this list of conditions and the following disclaimer.
>+ * 2. Redistributions in binary form must reproduce the above copyright
>+ *    notice, this list of conditions and the following disclaimer in the
>+ *    documentation and/or other materials provided with the distribution.
>+ *
>+ * THIS SOFTWARE IS PROVIDED BY APPLE INC. AND ITS CONTRIBUTORS ``AS IS''
>+ * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO,
>+ * THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
>+ * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL APPLE INC. OR ITS CONTRIBUTORS
>+ * BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
>+ * CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
>+ * SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
>+ * INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
>+ * CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
>+ * ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
>+ * THE POSSIBILITY OF SUCH DAMAGE.
>+ */
>+
>+#include "config.h"
>+#include "WHLSLEntryPointScaffolding.h"
>+
>+#if ENABLE(WEBGPU)
>+
>+#include "WHLSLGatherEntryPointItems.h"
>+
>+namespace WebCore {
>+
>+namespace WHLSL {
>+
>+namespace Metal {
>+
>+EntryPointScaffolding::EntryPointScaffolding(AST::FunctionDefinition& functionDefinition, Intrinsics& intrinsics)
>+    : m_functionDefinition(&functionDefinition)
>+    , m_intrinsics(&intrinsics)
>+{
>+    // FIXME: Implement this.
>+    gatherEntryPointItems(*m_intrinsics, *m_functionDefinition);
>+}
>+
>+String EntryPointScaffolding::helperTypes()
>+{
>+    // FIXME: Implement this.
>+    return String();
>+}
>+
>+String EntryPointScaffolding::signature()
>+{
>+    // FIXME: Implement this.
>+    return String();
>+}
>+
>+String EntryPointScaffolding::unpack()
>+{
>+    // FIXME: Implement this.
>+    return String();
>+}
>+
>+String EntryPointScaffolding::pack(const String&)
>+{
>+    // FIXME: Implement this.
>+    return String();
>+}
>+
>+}
>+
>+}
>+
>+}
>+
>+#endif
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.h b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.h
>new file mode 100644
>index 00000000000..1aa2da3bc4e
>--- /dev/null
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.h
>@@ -0,0 +1,65 @@
>+/*
>+ * Copyright (C) 2019 Apple Inc. All rights reserved.
>+ *
>+ * Redistribution and use in source and binary forms, with or without
>+ * modification, are permitted provided that the following conditions
>+ * are met:
>+ * 1. Redistributions of source code must retain the above copyright
>+ *    notice, this list of conditions and the following disclaimer.
>+ * 2. Redistributions in binary form must reproduce the above copyright
>+ *    notice, this list of conditions and the following disclaimer in the
>+ *    documentation and/or other materials provided with the distribution.
>+ *
>+ * THIS SOFTWARE IS PROVIDED BY APPLE INC. AND ITS CONTRIBUTORS ``AS IS''
>+ * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO,
>+ * THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
>+ * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL APPLE INC. OR ITS CONTRIBUTORS
>+ * BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
>+ * CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
>+ * SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
>+ * INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
>+ * CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
>+ * ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
>+ * THE POSSIBILITY OF SUCH DAMAGE.
>+ */
>+
>+#pragma once
>+
>+#if ENABLE(WEBGPU)
>+
>+namespace WebCore {
>+
>+namespace WHLSL {
>+
>+namespace AST {
>+
>+class FunctionDefinition;
>+
>+}
>+
>+class Intrinsics;
>+
>+namespace Metal {
>+
>+// FIXME: This needs to know about the pipeline state object to emit function prologues and epilogues.
>+class EntryPointScaffolding {
>+public:
>+    EntryPointScaffolding(AST::FunctionDefinition&, Intrinsics&);
>+
>+    String helperTypes();
>+    String signature();
>+    String unpack();
>+    String pack(const String& variableName);
>+
>+private:
>+    AST::FunctionDefinition* m_functionDefinition;
>+    Intrinsics* m_intrinsics;
>+};
>+
>+}
>+
>+}
>+
>+}
>+
>+#endif
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.cpp b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.cpp
>new file mode 100644
>index 00000000000..5cb79ae4464
>--- /dev/null
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.cpp
>@@ -0,0 +1,628 @@
>+/*
>+ * Copyright (C) 2019 Apple Inc. All rights reserved.
>+ *
>+ * Redistribution and use in source and binary forms, with or without
>+ * modification, are permitted provided that the following conditions
>+ * are met:
>+ * 1. Redistributions of source code must retain the above copyright
>+ *    notice, this list of conditions and the following disclaimer.
>+ * 2. Redistributions in binary form must reproduce the above copyright
>+ *    notice, this list of conditions and the following disclaimer in the
>+ *    documentation and/or other materials provided with the distribution.
>+ *
>+ * THIS SOFTWARE IS PROVIDED BY APPLE INC. AND ITS CONTRIBUTORS ``AS IS''
>+ * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO,
>+ * THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
>+ * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL APPLE INC. OR ITS CONTRIBUTORS
>+ * BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
>+ * CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
>+ * SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
>+ * INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
>+ * CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
>+ * ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
>+ * THE POSSIBILITY OF SUCH DAMAGE.
>+ */
>+
>+#include "config.h"
>+#include "WHLSLFunctionWriter.h"
>+
>+#if ENABLE(WEBGPU)
>+
>+#include "WHLSLArrayReferenceType.h"
>+#include "WHLSLArrayType.h"
>+#include "WHLSLAssignmentExpression.h"
>+#include "WHLSLBooleanLiteral.h"
>+#include "WHLSLBuiltInSemantic.h"
>+#include "WHLSLCallExpression.h"
>+#include "WHLSLCommaExpression.h"
>+#include "WHLSLDereferenceExpression.h"
>+#include "WHLSLDoWhileLoop.h"
>+#include "WHLSLEffectfulExpressionStatement.h"
>+#include "WHLSLEntryPointScaffolding.h"
>+#include "WHLSLEntryPointType.h"
>+#include "WHLSLFloatLiteral.h"
>+#include "WHLSLForLoop.h"
>+#include "WHLSLFunctionDeclaration.h"
>+#include "WHLSLFunctionDefinition.h"
>+#include "WHLSLIfStatement.h"
>+#include "WHLSLIntegerLiteral.h"
>+#include "WHLSLLogicalExpression.h"
>+#include "WHLSLLogicalNotExpression.h"
>+#include "WHLSLMakeArrayReferenceExpression.h"
>+#include "WHLSLMakePointerExpression.h"
>+#include "WHLSLNativeFunctionDeclaration.h"
>+#include "WHLSLNativeFunctionWriter.h"
>+#include "WHLSLNativeTypeDeclaration.h"
>+#include "WHLSLPointerType.h"
>+#include "WHLSLProgram.h"
>+#include "WHLSLReturn.h"
>+#include "WHLSLSwitchCase.h"
>+#include "WHLSLSwitchStatement.h"
>+#include "WHLSLTernaryExpression.h"
>+#include "WHLSLTypeNamer.h"
>+#include "WHLSLUnsignedIntegerLiteral.h"
>+#include "WHLSLVariableDeclaration.h"
>+#include "WHLSLVariableDeclarationsStatement.h"
>+#include "WHLSLVariableReference.h"
>+#include "WHLSLVisitor.h"
>+#include "WHLSLWhileLoop.h"
>+#include <wtf/HashMap.h>
>+#include <wtf/text/StringBuilder.h>
>+
>+namespace WebCore {
>+
>+namespace WHLSL {
>+
>+namespace Metal {
>+
>+class FunctionDeclarationWriter : public Visitor {
>+public:
>+    FunctionDeclarationWriter(TypeNamer& typeNamer, HashMap<AST::FunctionDeclaration*, String>& functionMapping)
>+        : m_typeNamer(typeNamer)
>+        , m_functionMapping(functionMapping)
>+    {
>+    }
>+
>+    virtual ~FunctionDeclarationWriter() = default;
>+
>+    String toString() { return m_stringBuilder.toString(); }
>+    void visit(AST::FunctionDeclaration& functionDeclaration) override
>+    {
>+        if (functionDeclaration.entryPointType())
>+            return;
>+
>+        auto iterator = m_functionMapping.find(&functionDeclaration);
>+        ASSERT(iterator != m_functionMapping.end());
>+        m_stringBuilder.append(makeString(m_typeNamer.mangledNameForType(functionDeclaration.type()), ' ', iterator->value, '('));
>+        for (size_t i = 0; i < functionDeclaration.parameters().size(); ++i) {
>+            if (i)
>+                m_stringBuilder.append(", ");
>+            m_stringBuilder.append(m_typeNamer.mangledNameForType(*functionDeclaration.parameters()[i].type()));
>+        }
>+        m_stringBuilder.append(");\n");
>+    }
>+
>+private:
>+    TypeNamer& m_typeNamer;
>+    HashMap<AST::FunctionDeclaration*, String>& m_functionMapping;
>+    StringBuilder m_stringBuilder;
>+};
>+
>+class FunctionDefinitionWriter : public Visitor {
>+public:
>+    FunctionDefinitionWriter(Intrinsics& intrinsics, TypeNamer& typeNamer, HashMap<AST::FunctionDeclaration*, String>& functionMapping)
>+        : m_intrinsics(intrinsics)
>+        , m_typeNamer(typeNamer)
>+        , m_functionMapping(functionMapping)
>+    {
>+    }
>+
>+    virtual ~FunctionDefinitionWriter() = default;
>+
>+    String toString() { return m_stringBuilder.toString(); }
>+
>+    void visit(AST::NativeFunctionDeclaration& nativeFunctionDeclaration) override
>+    {
>+        auto iterator = m_functionMapping.find(&nativeFunctionDeclaration);
>+        ASSERT(iterator != m_functionMapping.end());
>+        m_stringBuilder.append(writeNativeFunction(nativeFunctionDeclaration, iterator->value, m_typeNamer));
>+    }
>+
>+    void visit(AST::FunctionDefinition& functionDefinition) override
>+    {
>+        auto iterator = m_functionMapping.find(&functionDefinition);
>+        ASSERT(iterator != m_functionMapping.end());
>+        if (functionDefinition.entryPointType()) {
>+            m_entryPointScaffolding = EntryPointScaffolding(functionDefinition, m_intrinsics);
>+            m_stringBuilder.append(m_entryPointScaffolding->helperTypes());
>+            m_stringBuilder.append('\n');
>+            m_stringBuilder.append(makeString(m_entryPointScaffolding->signature(), " {"));
>+            m_stringBuilder.append(m_entryPointScaffolding->unpack());
>+            checkErrorAndVisit(functionDefinition.block());
>+            m_stringBuilder.append("}\n");
>+        } else {
>+            m_entryPointScaffolding = WTF::nullopt;
>+            m_stringBuilder.append(makeString(m_typeNamer.mangledNameForType(functionDefinition.type()), ' ', iterator->value, '('));
>+            for (size_t i = 0; i < functionDefinition.parameters().size(); ++i) {
>+                auto& parameter = functionDefinition.parameters()[i];
>+                if (i)
>+                    m_stringBuilder.append(", ");
>+                auto parameterName = generateNextVariableName();
>+                auto addResult = m_variableMapping.add(&parameter, parameterName);
>+                ASSERT_UNUSED(addResult, addResult.isNewEntry);
>+                m_stringBuilder.append(makeString(m_typeNamer.mangledNameForType(*parameter.type()), ' ', parameterName));
>+            }
>+            m_stringBuilder.append(") {\n");
>+            checkErrorAndVisit(functionDefinition.block());
>+            m_stringBuilder.append("}\n");
>+        }
>+    }
>+
>+private:
>+    void visit(AST::FunctionDeclaration&) override
>+    {
>+        ASSERT_NOT_REACHED();
>+    }
>+
>+    void visit(AST::Statement& statement) override
>+    {
>+        Visitor::visit(statement);
>+    }
>+
>+    void visit(AST::Block& block) override
>+    {
>+        m_stringBuilder.append("{\n");
>+        for (auto& statement : block.statements())
>+            checkErrorAndVisit(statement);
>+        m_stringBuilder.append("}\n");
>+    }
>+
>+    void visit(AST::Break&) override
>+    {
>+        m_stringBuilder.append("break;\n");
>+    }
>+
>+    void visit(AST::Continue&) override
>+    {
>+        // FIXME: Figure out which loop we're in, and run the increment code
>+        CRASH();
>+    }
>+
>+    void visit(AST::DoWhileLoop& doWhileLoop) override
>+    {
>+        m_stringBuilder.append("do {\n");
>+        checkErrorAndVisit(doWhileLoop.body());
>+        checkErrorAndVisit(doWhileLoop.conditional());
>+        m_stringBuilder.append(makeString("if (!", m_stack.takeLast(), ") break;\n"));
>+        m_stringBuilder.append(makeString("} while(true);\n"));
>+    }
>+
>+    void visit(AST::EffectfulExpressionStatement& effectfulExpressionStatement) override
>+    {
>+        checkErrorAndVisit(effectfulExpressionStatement.effectfulExpression());
>+        m_stack.takeLast(); // The statement is already effectful, so we don't need to do anything with the result.
>+    }
>+
>+    void visit(AST::Fallthrough&) override
>+    {
>+        m_stringBuilder.append("[[clang::fallthrough]];\n"); // FIXME: Make sure this is okay. Alternatively, we could do thing and just return here instead.
>+    }
>+
>+    void visit(AST::ForLoop& forLoop) override
>+    {
>+        WTF::visit(WTF::makeVisitor([&](AST::VariableDeclarationsStatement& variableDeclarationsStatement) {
>+            checkErrorAndVisit(variableDeclarationsStatement);
>+        }, [&](UniqueRef<AST::Expression>& expression) {
>+            checkErrorAndVisit(expression);
>+            m_stack.takeLast(); // We don't need to do anything with the result.
>+        }), forLoop.initialization());
>+
>+        m_stringBuilder.append("for ( ; ; ) {\n");
>+        if (forLoop.condition()) {
>+            checkErrorAndVisit(*forLoop.condition());
>+            m_stringBuilder.append(makeString("if (!", m_stack.takeLast(), ") break;\n"));
>+        }
>+        checkErrorAndVisit(forLoop.body());
>+        if (forLoop.increment())
>+            checkErrorAndVisit(*forLoop.increment());
>+        m_stringBuilder.append("}\n");
>+    }
>+
>+    void visit(AST::IfStatement& ifStatement) override
>+    {
>+        checkErrorAndVisit(ifStatement.conditional());
>+        m_stringBuilder.append(makeString("if (", m_stack.takeLast(), ") {\n"));
>+        checkErrorAndVisit(ifStatement.body());
>+        if (ifStatement.elseBody()) {
>+            m_stringBuilder.append("} else {\n");
>+            checkErrorAndVisit(*ifStatement.elseBody());
>+        }
>+        m_stringBuilder.append("}\n");
>+    }
>+
>+    void visit(AST::Return& returnStatement) override
>+    {
>+        if (returnStatement.value()) {
>+            checkErrorAndVisit(*returnStatement.value());
>+            if (m_entryPointScaffolding) {
>+                auto variableName = generateNextVariableName();
>+                m_stringBuilder.append(m_entryPointScaffolding->pack(variableName));
>+                m_stringBuilder.append(makeString("return ", variableName, ";\n"));
>+            } else
>+                m_stringBuilder.append(makeString("return ", m_stack.takeLast(), ";\n"));
>+        } else
>+            m_stringBuilder.append("return;\n");
>+    }
>+
>+    void visit(AST::SwitchStatement& switchStatement) override
>+    {
>+        checkErrorAndVisit(switchStatement.value());
>+
>+        m_stringBuilder.append(makeString("switch (", m_stack.takeLast(), ") {"));
>+        for (auto& switchCase : switchStatement.switchCases())
>+            checkErrorAndVisit(switchCase);
>+        m_stringBuilder.append("}\n");
>+    }
>+
>+    void visit(AST::SwitchCase& switchCase) override
>+    {
>+        if (switchCase.value())
>+            m_stringBuilder.append(makeString("case ", constantExpressionString(*switchCase.value()), ":\n"));
>+        else
>+            m_stringBuilder.append("default:\n");
>+        // FIXME: Figure out whether we need to break or fallthrough.
>+        CRASH();
>+        checkErrorAndVisit(switchCase.block());
>+    }
>+
>+    void visit(AST::Trap&) override
>+    {
>+        // FIXME: Implement this
>+        CRASH();
>+    }
>+
>+    void visit(AST::VariableDeclarationsStatement& variableDeclarationsStatement) override
>+    {
>+        Visitor::visit(variableDeclarationsStatement);
>+    }
>+
>+    void visit(AST::WhileLoop& whileLoop) override
>+    {
>+        m_stringBuilder.append(makeString("while (true) {\n"));
>+        checkErrorAndVisit(whileLoop.conditional());
>+        m_stringBuilder.append(makeString("if (!", m_stack.takeLast(), ") break;\n"));
>+        checkErrorAndVisit(whileLoop.body());
>+        m_stringBuilder.append("}\n");
>+    }
>+
>+    void visit(AST::IntegerLiteral& integerLiteral) override
>+    {
>+        ASSERT(integerLiteral.resolvedType());
>+        auto variableName = generateNextVariableName();
>+        auto mangledTypeName = m_typeNamer.mangledNameForType(*integerLiteral.resolvedType());
>+        m_stringBuilder.append(makeString(mangledTypeName, ' ', variableName, " = static_cast<", mangledTypeName, ">(", integerLiteral.value(), ");\n"));
>+        m_stack.append(variableName);
>+    }
>+
>+    void visit(AST::UnsignedIntegerLiteral& unsignedIntegerLiteral) override
>+    {
>+        ASSERT(unsignedIntegerLiteral.resolvedType());
>+        auto variableName = generateNextVariableName();
>+        auto mangledTypeName = m_typeNamer.mangledNameForType(*unsignedIntegerLiteral.resolvedType());
>+        m_stringBuilder.append(makeString(mangledTypeName, ' ', variableName, " = static_cast<", mangledTypeName, ">(", unsignedIntegerLiteral.value(), ");\n"));
>+        m_stack.append(variableName);
>+    }
>+
>+    void visit(AST::FloatLiteral& floatLiteral) override
>+    {
>+        ASSERT(floatLiteral.resolvedType());
>+        auto variableName = generateNextVariableName();
>+        auto mangledTypeName = m_typeNamer.mangledNameForType(*floatLiteral.resolvedType());
>+        m_stringBuilder.append(makeString(mangledTypeName, ' ', variableName, " = static_cast<", mangledTypeName, ">(", floatLiteral.value(), ");\n"));
>+        m_stack.append(variableName);
>+    }
>+
>+    void visit(AST::NullLiteral& nullLiteral) override
>+    {
>+        ASSERT(nullLiteral.resolvedType());
>+        auto& unifyNode = nullLiteral.resolvedType()->unifyNode();
>+        ASSERT(is<AST::UnnamedType>(unifyNode));
>+        auto& unnamedType = downcast<AST::UnnamedType>(unifyNode);
>+        bool isArrayReferenceType = is<AST::ArrayReferenceType>(unnamedType);
>+
>+        auto variableName = generateNextVariableName();
>+        m_stringBuilder.append(makeString(m_typeNamer.mangledNameForType(*nullLiteral.resolvedType()), ' ', variableName, " = "));
>+        if (isArrayReferenceType)
>+            m_stringBuilder.append("{ nullptr, 0 }");
>+        else
>+            m_stringBuilder.append("nullptr");
>+        m_stringBuilder.append(";\n");
>+        m_stack.append(variableName);
>+    }
>+
>+    void visit(AST::BooleanLiteral& booleanLiteral) override
>+    {
>+        ASSERT(booleanLiteral.resolvedType());
>+        auto variableName = generateNextVariableName();
>+        auto mangledTypeName = m_typeNamer.mangledNameForType(*booleanLiteral.resolvedType());
>+        m_stringBuilder.append(makeString(mangledTypeName, ' ', variableName, " = static_cast<", mangledTypeName, ">(", booleanLiteral.value() ? "true" : "false", ");\n"));
>+        m_stack.append(variableName);
>+    }
>+
>+    void visit(AST::EnumerationMemberLiteral& enumerationMemberLiteral) override
>+    {
>+        ASSERT(enumerationMemberLiteral.resolvedType());
>+        ASSERT(enumerationMemberLiteral.enumerationDefinition());
>+        ASSERT(enumerationMemberLiteral.enumerationDefinition());
>+        auto variableName = generateNextVariableName();
>+        auto mangledTypeName = m_typeNamer.mangledNameForType(*enumerationMemberLiteral.resolvedType());
>+        m_stringBuilder.append(makeString(mangledTypeName, ' ', variableName, " = ", mangledTypeName, '.', m_typeNamer.mangledNameForEnumerationMember(*enumerationMemberLiteral.enumerationMember()), ";\n"));
>+        m_stack.append(variableName);
>+    }
>+
>+    void visit(AST::Expression& expression) override
>+    {
>+        Visitor::visit(expression);
>+    }
>+
>+    void visit(AST::DotExpression&) override
>+    {
>+        // This should be lowered already.
>+        ASSERT_NOT_REACHED();
>+    }
>+
>+    void visit(AST::IndexExpression&) override
>+    {
>+        // This should be lowered already.
>+        ASSERT_NOT_REACHED();
>+    }
>+
>+    void visit(AST::PropertyAccessExpression&) override
>+    {
>+        ASSERT_NOT_REACHED();
>+    }
>+
>+    void visit(AST::VariableDeclaration& variableDeclaration) override
>+    {
>+        ASSERT(variableDeclaration.type());
>+        if (variableDeclaration.initializer())
>+            checkErrorAndVisit(*variableDeclaration.initializer());
>+        else {
>+            // FIXME: Zero-fill the variable.
>+            CRASH();
>+        }
>+        // FIXME: Implement qualifiers.
>+        auto variableName = generateNextVariableName();
>+        auto addResult = m_variableMapping.add(&variableDeclaration, variableName);
>+        ASSERT_UNUSED(addResult, addResult.isNewEntry);
>+        m_stringBuilder.append(makeString(m_typeNamer.mangledNameForType(*variableDeclaration.type()), ' ', variableName, " = ", m_stack.takeLast(), ";\n"));
>+        m_stack.append(variableName);
>+    }
>+
>+    void visit(AST::AssignmentExpression& assignmentExpression) override
>+    {
>+        checkErrorAndVisit(assignmentExpression.left());
>+        auto leftName = m_stack.takeLast();
>+        checkErrorAndVisit(assignmentExpression.right());
>+        auto rightName = m_stack.takeLast();
>+        m_stringBuilder.append(makeString(leftName, " = ", rightName, ";\n"));
>+    }
>+
>+    void visit(AST::CallExpression& callExpression) override
>+    {
>+        Vector<String> argumentNames;
>+        for (auto& argument : callExpression.arguments()) {
>+            checkErrorAndVisit(argument);
>+            argumentNames.append(m_stack.takeLast());
>+        }
>+        ASSERT(callExpression.resolvedType());
>+        ASSERT(callExpression.function());
>+        auto iterator = m_functionMapping.find(callExpression.function());
>+        ASSERT(iterator != m_functionMapping.end());
>+        auto variableName = generateNextVariableName();
>+        m_stringBuilder.append(makeString(m_typeNamer.mangledNameForType(*callExpression.resolvedType()), ' ', variableName, " = ", iterator->value, '('));
>+        for (size_t i = 0; i < argumentNames.size(); ++i) {
>+            if (i)
>+                m_stringBuilder.append(", ");
>+            m_stringBuilder.append(argumentNames[i]);
>+        }
>+        m_stringBuilder.append(");\n");
>+        m_stack.append(variableName);
>+    }
>+
>+    void visit(AST::CommaExpression& commaExpression) override
>+    {
>+        String result;
>+        for (auto& expression : commaExpression.list()) {
>+            checkErrorAndVisit(expression);
>+            result = m_stack.takeLast();
>+        }
>+        m_stack.append(result);
>+    }
>+
>+    void visit(AST::DereferenceExpression& dereferenceExpression) override
>+    {
>+        checkErrorAndVisit(dereferenceExpression.pointer());
>+        auto right = m_stack.takeLast();
>+        ASSERT(dereferenceExpression.resolvedType());
>+        auto variableName = generateNextVariableName();
>+        m_stringBuilder.append(makeString(m_typeNamer.mangledNameForType(*dereferenceExpression.resolvedType()), ' ', variableName, " = *", right, ";\n"));
>+        m_stack.append(variableName);
>+    }
>+
>+    void visit(AST::LogicalExpression& logicalExpression) override
>+    {
>+        checkErrorAndVisit(logicalExpression.left());
>+        auto left = m_stack.takeLast();
>+        checkErrorAndVisit(logicalExpression.right());
>+        auto right = m_stack.takeLast();
>+        ASSERT(logicalExpression.resolvedType());
>+        auto variableName = generateNextVariableName();
>+        m_stringBuilder.append(makeString(m_typeNamer.mangledNameForType(*logicalExpression.resolvedType()), ' ', variableName, " = ", left, ' '));
>+        switch (logicalExpression.type()) {
>+        case AST::LogicalExpression::Type::And:
>+            m_stringBuilder.append("&&");
>+            break;
>+        default:
>+            ASSERT(logicalExpression.type() == AST::LogicalExpression::Type::Or);
>+            m_stringBuilder.append("||");
>+            break;
>+        }
>+        m_stringBuilder.append(makeString(' ', right, ";\n"));
>+        m_stack.append(variableName);
>+    }
>+
>+    void visit(AST::LogicalNotExpression& logicalNotExpression) override
>+    {
>+        checkErrorAndVisit(logicalNotExpression.operand());
>+        auto operand = m_stack.takeLast();
>+        ASSERT(logicalNotExpression.resolvedType());
>+        auto variableName = generateNextVariableName();
>+        m_stringBuilder.append(makeString(m_typeNamer.mangledNameForType(*logicalNotExpression.resolvedType()), ' ', variableName, " = !", operand, ";\n"));
>+        m_stack.append(variableName);
>+    }
>+
>+    void visit(AST::MakeArrayReferenceExpression& makeArrayReferenceExpression) override
>+    {
>+        checkErrorAndVisit(makeArrayReferenceExpression.lValue());
>+        auto lValue = m_stack.takeLast();
>+        ASSERT(makeArrayReferenceExpression.resolvedType());
>+        auto variableName = generateNextVariableName();
>+        auto mangledTypeName = m_typeNamer.mangledNameForType(*makeArrayReferenceExpression.resolvedType());
>+        if (is<AST::PointerType>(*makeArrayReferenceExpression.resolvedType()))
>+            m_stringBuilder.append(makeString(mangledTypeName, ' ', variableName, " = { ", lValue, ", 1 };\n"));
>+        else if (is<AST::ArrayType>(*makeArrayReferenceExpression.resolvedType())) {
>+            auto& arrayType = downcast<AST::ArrayType>(*makeArrayReferenceExpression.resolvedType());
>+            m_stringBuilder.append(makeString(mangledTypeName, ' ', variableName, " = { &(", lValue, "[0]), ", arrayType.numElements(), " };\n"));
>+        } else
>+            m_stringBuilder.append(makeString(mangledTypeName, ' ', variableName, " = { &", lValue, ", 1 };\n"));
>+        m_stack.append(variableName);
>+    }
>+
>+    void visit(AST::MakePointerExpression& makePointerExpression) override
>+    {
>+        checkErrorAndVisit(makePointerExpression.lValue());
>+        auto lValue = m_stack.takeLast();
>+        ASSERT(makePointerExpression.resolvedType());
>+        auto variableName = generateNextVariableName();
>+        m_stringBuilder.append(makeString(m_typeNamer.mangledNameForType(*makePointerExpression.resolvedType()), ' ', variableName, " = &", lValue, ";\n"));
>+        m_stack.append(variableName);
>+    }
>+
>+    void visit(AST::ReadModifyWriteExpression&) override
>+    {
>+        // This should be lowered already.
>+        ASSERT_NOT_REACHED();
>+    }
>+
>+    void visit(AST::TernaryExpression& ternaryExpression) override
>+    {
>+        checkErrorAndVisit(ternaryExpression.predicate());
>+        auto check = m_stack.takeLast();
>+
>+        ASSERT(ternaryExpression.resolvedType());
>+        auto variableName = generateNextVariableName();
>+        m_stringBuilder.append(makeString(m_typeNamer.mangledNameForType(*ternaryExpression.resolvedType()), ' ', variableName, ";\n"));
>+
>+        m_stringBuilder.append(makeString("if (", check, ") {\n"));
>+        checkErrorAndVisit(ternaryExpression.bodyExpression());
>+        m_stringBuilder.append(makeString(variableName, " = ", m_stack.takeLast(), ";\n"));
>+        m_stringBuilder.append("} else {\n");
>+        checkErrorAndVisit(ternaryExpression.elseExpression());
>+        m_stringBuilder.append(makeString(variableName, " = ", m_stack.takeLast(), ";\n"));
>+        m_stringBuilder.append("}\n");
>+        m_stack.append(variableName);
>+    }
>+
>+    void visit(AST::VariableReference& variableReference) override
>+    {
>+        ASSERT(variableReference.variable());
>+        auto iterator = m_variableMapping.find(variableReference.variable());
>+        ASSERT(iterator != m_variableMapping.end());
>+        m_stack.append(iterator->value);
>+    }
>+
>+    String constantExpressionString(AST::ConstantExpression& constantExpression)
>+    {
>+        String result;
>+        constantExpression.visit(WTF::makeVisitor([&](AST::IntegerLiteral& integerLiteral) {
>+            result = makeString("", integerLiteral.value());
>+        }, [&](AST::UnsignedIntegerLiteral& unsignedIntegerLiteral) {
>+            result = makeString("", unsignedIntegerLiteral.value());
>+        }, [&](AST::FloatLiteral& floatLiteral) {
>+            result = makeString("", floatLiteral.value());
>+        }, [&](AST::NullLiteral&) {
>+            result = "nullptr"_str;
>+        }, [&](AST::BooleanLiteral& booleanLiteral) {
>+            result = booleanLiteral.value() ? "true"_str : "false"_str;
>+        }, [&](AST::EnumerationMemberLiteral& enumerationMemberLiteral) {
>+            ASSERT(enumerationMemberLiteral.enumerationDefinition());
>+            ASSERT(enumerationMemberLiteral.enumerationDefinition());
>+            result = makeString(m_typeNamer.mangledNameForType(*enumerationMemberLiteral.enumerationDefinition()), '.', m_typeNamer.mangledNameForEnumerationMember(*enumerationMemberLiteral.enumerationMember()));
>+        }));
>+        return result;
>+    }
>+
>+    String generateNextVariableName()
>+    {
>+        return makeString("variable", m_variableCount++);
>+    }
>+
>+private:
>+    Intrinsics& m_intrinsics;
>+    TypeNamer& m_typeNamer;
>+    HashMap<AST::FunctionDeclaration*, String>& m_functionMapping;
>+    HashMap<AST::VariableDeclaration*, String> m_variableMapping;
>+    StringBuilder m_stringBuilder;
>+    Vector<String> m_stack;
>+    Optional<EntryPointScaffolding> m_entryPointScaffolding;
>+    unsigned m_variableCount { 0 };
>+};
>+
>+String metalFunctions(Program& program, TypeNamer& typeNamer)
>+{
>+    StringBuilder stringBuilder;
>+
>+    unsigned numFunctions = 0;
>+    HashMap<AST::FunctionDeclaration*, String> functionMapping;
>+    for (auto& nativeFunctionDeclaration : program.nativeFunctionDeclarations()) {
>+        auto addResult = functionMapping.add(&nativeFunctionDeclaration, makeString("function", numFunctions++));
>+        ASSERT_UNUSED(addResult, addResult.isNewEntry);
>+    }
>+    for (auto& functionDefinition : program.functionDefinitions()) {
>+        auto addResult = functionMapping.add(&functionDefinition, makeString("function", numFunctions++));
>+        ASSERT_UNUSED(addResult, addResult.isNewEntry);
>+    }
>+
>+    {
>+        FunctionDeclarationWriter functionDeclarationWriter(typeNamer, functionMapping);
>+        for (auto& nativeFunctionDeclaration : program.nativeFunctionDeclarations())
>+            functionDeclarationWriter.visit(nativeFunctionDeclaration);
>+        for (auto& functionDefinition : program.functionDefinitions()) {
>+            if (!functionDefinition->entryPointType())
>+                functionDeclarationWriter.visit(functionDefinition);
>+        }
>+        stringBuilder.append(functionDeclarationWriter.toString());
>+    }
>+
>+    stringBuilder.append('\n');
>+
>+    {
>+        FunctionDefinitionWriter functionDefinitionWriter(program.intrinsics(), typeNamer, functionMapping);
>+        for (auto& nativeFunctionDeclaration : program.nativeFunctionDeclarations())
>+            functionDefinitionWriter.visit(nativeFunctionDeclaration);
>+        for (auto& functionDefinition : program.functionDefinitions())
>+            functionDefinitionWriter.visit(functionDefinition);
>+        stringBuilder.append(functionDefinitionWriter.toString());
>+    }
>+
>+    return stringBuilder.toString();
>+}
>+
>+} // namespace Metal
>+
>+} // namespace WHLSL
>+
>+} // namespace WebCore
>+
>+#endif
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.h b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.h
>new file mode 100644
>index 00000000000..9960e5e48d5
>--- /dev/null
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.h
>@@ -0,0 +1,48 @@
>+/*
>+ * Copyright (C) 2019 Apple Inc. All rights reserved.
>+ *
>+ * Redistribution and use in source and binary forms, with or without
>+ * modification, are permitted provided that the following conditions
>+ * are met:
>+ * 1. Redistributions of source code must retain the above copyright
>+ *    notice, this list of conditions and the following disclaimer.
>+ * 2. Redistributions in binary form must reproduce the above copyright
>+ *    notice, this list of conditions and the following disclaimer in the
>+ *    documentation and/or other materials provided with the distribution.
>+ *
>+ * THIS SOFTWARE IS PROVIDED BY APPLE INC. AND ITS CONTRIBUTORS ``AS IS''
>+ * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO,
>+ * THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
>+ * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL APPLE INC. OR ITS CONTRIBUTORS
>+ * BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
>+ * CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
>+ * SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
>+ * INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
>+ * CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
>+ * ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
>+ * THE POSSIBILITY OF SUCH DAMAGE.
>+ */
>+
>+#pragma once
>+
>+#if ENABLE(WEBGPU)
>+
>+namespace WebCore {
>+
>+namespace WHLSL {
>+
>+class Program;
>+
>+namespace Metal {
>+
>+class TypeNamer;
>+
>+String metalFunctions(Program&, TypeNamer&);
>+
>+}
>+
>+}
>+
>+}
>+
>+#endif
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.cpp b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.cpp
>new file mode 100644
>index 00000000000..0277be5ddae
>--- /dev/null
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.cpp
>@@ -0,0 +1,68 @@
>+/*
>+ * Copyright (C) 2019 Apple Inc. All rights reserved.
>+ *
>+ * Redistribution and use in source and binary forms, with or without
>+ * modification, are permitted provided that the following conditions
>+ * are met:
>+ * 1. Redistributions of source code must retain the above copyright
>+ *    notice, this list of conditions and the following disclaimer.
>+ * 2. Redistributions in binary form must reproduce the above copyright
>+ *    notice, this list of conditions and the following disclaimer in the
>+ *    documentation and/or other materials provided with the distribution.
>+ *
>+ * THIS SOFTWARE IS PROVIDED BY APPLE INC. AND ITS CONTRIBUTORS ``AS IS''
>+ * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO,
>+ * THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
>+ * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL APPLE INC. OR ITS CONTRIBUTORS
>+ * BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
>+ * CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
>+ * SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
>+ * INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
>+ * CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
>+ * ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
>+ * THE POSSIBILITY OF SUCH DAMAGE.
>+ */
>+
>+#include "config.h"
>+#include "WHLSLMetalCodeGenerator.h"
>+
>+#if ENABLE(WEBGPU)
>+
>+#include "WHLSLFunctionWriter.h"
>+#include "WHLSLTypeNamer.h"
>+#include <wtf/text/StringBuilder.h>
>+
>+namespace WebCore {
>+
>+namespace WHLSL {
>+
>+namespace Metal {
>+
>+String generateMetalCode(Program& program)
>+{
>+    StringBuilder stringBuilder;
>+    stringBuilder.append("// This was generated by WebKit's WHLSL -> MSL code generator. Do not edit!\n");
>+    stringBuilder.append("\n");
>+    stringBuilder.append("#include <metal_stdlib>\n");
>+    stringBuilder.append("#include <metal_atomic>\n");
>+    stringBuilder.append("#include <metal_math>\n");
>+    stringBuilder.append("#include <metal_relational>\n");
>+    stringBuilder.append("#include <metal_compute>\n");
>+    stringBuilder.append("#include <metal_texture>\n");
>+    stringBuilder.append("\n");
>+    stringBuilder.append("using namespace metal;\n");
>+    stringBuilder.append("\n");
>+
>+    TypeNamer typeNamer(program);
>+    stringBuilder.append(typeNamer.metalTypes());
>+    stringBuilder.append(metalFunctions(program, typeNamer));
>+    return stringBuilder.toString();
>+}
>+
>+}
>+
>+}
>+
>+}
>+
>+#endif
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.h b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.h
>new file mode 100644
>index 00000000000..e088e010449
>--- /dev/null
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.h
>@@ -0,0 +1,49 @@
>+/*
>+ * Copyright (C) 2019 Apple Inc. All rights reserved.
>+ *
>+ * Redistribution and use in source and binary forms, with or without
>+ * modification, are permitted provided that the following conditions
>+ * are met:
>+ * 1. Redistributions of source code must retain the above copyright
>+ *    notice, this list of conditions and the following disclaimer.
>+ * 2. Redistributions in binary form must reproduce the above copyright
>+ *    notice, this list of conditions and the following disclaimer in the
>+ *    documentation and/or other materials provided with the distribution.
>+ *
>+ * THIS SOFTWARE IS PROVIDED BY APPLE INC. AND ITS CONTRIBUTORS ``AS IS''
>+ * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO,
>+ * THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
>+ * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL APPLE INC. OR ITS CONTRIBUTORS
>+ * BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
>+ * CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
>+ * SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
>+ * INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
>+ * CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
>+ * ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
>+ * THE POSSIBILITY OF SUCH DAMAGE.
>+ */
>+
>+#pragma once
>+
>+#if ENABLE(WEBGPU)
>+
>+#include <wtf/text/WTFString.h>
>+
>+namespace WebCore {
>+
>+namespace WHLSL {
>+
>+class Program;
>+
>+namespace Metal {
>+
>+// FIXME: This needs to know about the pipeline state object to emit function prologues and epilogues.
>+String generateMetalCode(Program&);
>+
>+}
>+
>+}
>+
>+}
>+
>+#endif
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLNativeFunctionWriter.cpp b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLNativeFunctionWriter.cpp
>new file mode 100644
>index 00000000000..bffec631b7a
>--- /dev/null
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLNativeFunctionWriter.cpp
>@@ -0,0 +1,394 @@
>+/*
>+ * Copyright (C) 2019 Apple Inc. All rights reserved.
>+ *
>+ * Redistribution and use in source and binary forms, with or without
>+ * modification, are permitted provided that the following conditions
>+ * are met:
>+ * 1. Redistributions of source code must retain the above copyright
>+ *    notice, this list of conditions and the following disclaimer.
>+ * 2. Redistributions in binary form must reproduce the above copyright
>+ *    notice, this list of conditions and the following disclaimer in the
>+ *    documentation and/or other materials provided with the distribution.
>+ *
>+ * THIS SOFTWARE IS PROVIDED BY APPLE INC. AND ITS CONTRIBUTORS ``AS IS''
>+ * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO,
>+ * THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
>+ * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL APPLE INC. OR ITS CONTRIBUTORS
>+ * BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
>+ * CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
>+ * SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
>+ * INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
>+ * CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
>+ * ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
>+ * THE POSSIBILITY OF SUCH DAMAGE.
>+ */
>+
>+#include "config.h"
>+#include "WHLSLNativeFunctionWriter.h"
>+
>+#if ENABLE(WEBGPU)
>+
>+#include "WHLSLAddressSpace.h"
>+#include "WHLSLNamedType.h"
>+#include "WHLSLNativeFunctionDeclaration.h"
>+#include "WHLSLNativeTypeDeclaration.h"
>+#include "WHLSLTypeNamer.h"
>+#include "WHLSLUnnamedType.h"
>+#include "WHLSLVariableDeclaration.h"
>+
>+namespace WebCore {
>+
>+namespace WHLSL {
>+
>+namespace Metal {
>+
>+static String getNativeName(AST::UnnamedType& unnamedType, TypeNamer& typeNamer)
>+{
>+    ASSERT(is<AST::NamedType>(unnamedType.unifyNode()));
>+    auto& namedType = downcast<AST::NamedType>(unnamedType.unifyNode());
>+    ASSERT(is<AST::NativeTypeDeclaration>(namedType));
>+    auto& nativeTypeDeclaration = downcast<AST::NativeTypeDeclaration>(namedType);
>+    return typeNamer.mangledNameForType(nativeTypeDeclaration);
>+}
>+
>+static String mapFunctionName(String& functionName)
>+{
>+    if (functionName == "ddx")
>+        return "dfdx"_str;
>+    if (functionName == "ddy")
>+        return "dfdy"_str;
>+    if (functionName == "asint")
>+        return "as_type<int32_t>"_str;
>+    if (functionName == "asuint")
>+        return "as_type<uint32_t>"_str;
>+    if (functionName == "asfloat")
>+        return "as_type<float>"_str;
>+    return functionName;
>+}
>+
>+static String convertAddressSpace(AST::AddressSpace addressSpace)
>+{
>+    switch (addressSpace) {
>+    case AST::AddressSpace::Constant:
>+        return "constant"_str;
>+    case AST::AddressSpace::Device:
>+        return "device"_str;
>+    case AST::AddressSpace::Threadgroup:
>+        return "threadgroup"_str;
>+    default:
>+        ASSERT(addressSpace == AST::AddressSpace::Thread);
>+        return "thread"_str;
>+    }
>+}
>+
>+String writeNativeFunction(AST::NativeFunctionDeclaration& nativeFunctionDeclaration, String& outputFunctionName, TypeNamer& typeNamer)
>+{
>+    StringBuilder stringBuilder;
>+    if (nativeFunctionDeclaration.isCast()) {
>+        ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
>+        auto metalParameterName = getNativeName(*nativeFunctionDeclaration.parameters()[0].type(), typeNamer);
>+        auto metalReturnName = getNativeName(nativeFunctionDeclaration.type(), typeNamer);
>+        if (metalParameterName != "atomic_int"_str && metalParameterName != "atomic_uint"_str) {
>+            stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameterName, "x) {\n"));
>+            stringBuilder.append(makeString("    return static_cast<", metalReturnName, ">(x);\n"));
>+            stringBuilder.append("}\n");
>+            return stringBuilder.toString();
>+        }
>+
>+        stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameterName, "x) {\n"));
>+        stringBuilder.append("    return atomic_load_explicit(&x, memory_order_relaxed);\n");
>+        stringBuilder.append("}\n");
>+        return stringBuilder.toString();
>+    }
>+
>+    if (nativeFunctionDeclaration.name().startsWith("operator."_str)) {
>+        if (nativeFunctionDeclaration.name().endsWith("="_str)) {
>+            ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
>+            auto metalParameter1Name = getNativeName(*nativeFunctionDeclaration.parameters()[0].type(), typeNamer);
>+            auto metalParameter2Name = getNativeName(*nativeFunctionDeclaration.parameters()[1].type(), typeNamer);
>+            auto metalReturnName = getNativeName(nativeFunctionDeclaration.type(), typeNamer);
>+            auto fieldName = nativeFunctionDeclaration.name().substring("operator."_str.length());
>+            fieldName = fieldName.substring(0, fieldName.length() - 1);
>+            stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, "v, ", metalParameter2Name, " n) {\n"));
>+            stringBuilder.append(makeString("    v.", fieldName, " = n;\n"));
>+            stringBuilder.append(makeString("    return v;\n"));
>+            stringBuilder.append("}\n");
>+            return stringBuilder.toString();
>+        }
>+
>+        ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
>+        auto metalParameterName = getNativeName(*nativeFunctionDeclaration.parameters()[0].type(), typeNamer);
>+        auto metalReturnName = getNativeName(nativeFunctionDeclaration.type(), typeNamer);
>+        stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameterName, "v) {\n"));
>+        stringBuilder.append(makeString("    return v.", nativeFunctionDeclaration.name().substring("operator."_str.length()), ";\n"));
>+        stringBuilder.append("}\n");
>+        return stringBuilder.toString();
>+
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "operator[]") {
>+        ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
>+        auto metalParameter1Name = getNativeName(*nativeFunctionDeclaration.parameters()[0].type(), typeNamer);
>+        auto metalParameter2Name = getNativeName(*nativeFunctionDeclaration.parameters()[1].type(), typeNamer);
>+        auto metalReturnName = getNativeName(nativeFunctionDeclaration.type(), typeNamer);
>+        stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, "m, ", metalParameter2Name, " i) {\n"));
>+        stringBuilder.append(makeString("    return m[i];\n"));
>+        stringBuilder.append("}\n");
>+        return stringBuilder.toString();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "operator[]=") {
>+        ASSERT(nativeFunctionDeclaration.parameters().size() == 3);
>+        auto metalParameter1Name = getNativeName(*nativeFunctionDeclaration.parameters()[0].type(), typeNamer);
>+        auto metalParameter2Name = getNativeName(*nativeFunctionDeclaration.parameters()[1].type(), typeNamer);
>+        auto metalParameter3Name = getNativeName(*nativeFunctionDeclaration.parameters()[2].type(), typeNamer);
>+        auto metalReturnName = getNativeName(nativeFunctionDeclaration.type(), typeNamer);
>+        stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, "m, ", metalParameter2Name, " i, ", metalParameter3Name, " v) {\n"));
>+        stringBuilder.append(makeString("    m[i] = v;\n"));
>+        stringBuilder.append(makeString("    return m;\n"));
>+        stringBuilder.append("}\n");
>+        return stringBuilder.toString();
>+    }
>+
>+    if (nativeFunctionDeclaration.isOperator()) {
>+        if (nativeFunctionDeclaration.parameters().size() == 1) {
>+            auto operatorName = nativeFunctionDeclaration.name().substring("operator"_str.length());
>+            auto metalParameterName = getNativeName(*nativeFunctionDeclaration.parameters()[0].type(), typeNamer);
>+            auto metalReturnName = getNativeName(nativeFunctionDeclaration.type(), typeNamer);
>+            stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameterName, "x) {\n"));
>+            stringBuilder.append(makeString("    return ", operatorName, "x;\n"));
>+            stringBuilder.append("}\n");
>+            return stringBuilder.toString();
>+        }
>+
>+        ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
>+        auto operatorName = nativeFunctionDeclaration.name().substring("operator"_str.length());
>+        auto metalParameter1Name = getNativeName(*nativeFunctionDeclaration.parameters()[0].type(), typeNamer);
>+        auto metalParameter2Name = getNativeName(*nativeFunctionDeclaration.parameters()[1].type(), typeNamer);
>+        auto metalReturnName = getNativeName(nativeFunctionDeclaration.type(), typeNamer);
>+        stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, "x, ", metalParameter2Name, " y) {\n"));
>+        stringBuilder.append(makeString("    return x ", operatorName, " y;\n"));
>+        stringBuilder.append("}\n");
>+        return stringBuilder.toString();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "cos"
>+        || nativeFunctionDeclaration.name() == "sin"
>+        || nativeFunctionDeclaration.name() == "tan"
>+        || nativeFunctionDeclaration.name() == "acos"
>+        || nativeFunctionDeclaration.name() == "asin"
>+        || nativeFunctionDeclaration.name() == "atan"
>+        || nativeFunctionDeclaration.name() == "cosh"
>+        || nativeFunctionDeclaration.name() == "sinh"
>+        || nativeFunctionDeclaration.name() == "tanh"
>+        || nativeFunctionDeclaration.name() == "ceil"
>+        || nativeFunctionDeclaration.name() == "exp"
>+        || nativeFunctionDeclaration.name() == "floor"
>+        || nativeFunctionDeclaration.name() == "log"
>+        || nativeFunctionDeclaration.name() == "round"
>+        || nativeFunctionDeclaration.name() == "trunc"
>+        || nativeFunctionDeclaration.name() == "ddx"
>+        || nativeFunctionDeclaration.name() == "ddy"
>+        || nativeFunctionDeclaration.name() == "isnormal"
>+        || nativeFunctionDeclaration.name() == "isfinite"
>+        || nativeFunctionDeclaration.name() == "isinf"
>+        || nativeFunctionDeclaration.name() == "isnan"
>+        || nativeFunctionDeclaration.name() == "asint"
>+        || nativeFunctionDeclaration.name() == "asuint"
>+        || nativeFunctionDeclaration.name() == "asfloat") {
>+        ASSERT(nativeFunctionDeclaration.parameters().size() == 1);
>+        auto metalParameterName = getNativeName(*nativeFunctionDeclaration.parameters()[0].type(), typeNamer);
>+        auto metalReturnName = getNativeName(nativeFunctionDeclaration.type(), typeNamer);
>+        stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameterName, "x) {\n"));
>+        stringBuilder.append(makeString("    return ", mapFunctionName(nativeFunctionDeclaration.name()), "(x);\n"));
>+        stringBuilder.append("}\n");
>+        return stringBuilder.toString();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "pow" || nativeFunctionDeclaration.name() == "atan2") {
>+        ASSERT(nativeFunctionDeclaration.parameters().size() == 2);
>+        auto metalParameter1Name = getNativeName(*nativeFunctionDeclaration.parameters()[0].type(), typeNamer);
>+        auto metalParameter2Name = getNativeName(*nativeFunctionDeclaration.parameters()[1].type(), typeNamer);
>+        auto metalReturnName = getNativeName(nativeFunctionDeclaration.type(), typeNamer);
>+        stringBuilder.append(makeString(metalReturnName, ' ', outputFunctionName, '(', metalParameter1Name, "x, ", metalParameter2Name, " y) {\n"));
>+        stringBuilder.append(makeString("    return ", nativeFunctionDeclaration.name(), "(x, y);\n"));
>+        stringBuilder.append("}\n");
>+        return stringBuilder.toString();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "f16tof32" || nativeFunctionDeclaration.name() == "f32tof16") {
>+        // FIXME: Implement this
>+        CRASH();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "AllMemoryBarrierWithGroupSync") {
>+        ASSERT(!nativeFunctionDeclaration.parameters().size());
>+        stringBuilder.append(makeString("void ", outputFunctionName, "() {\n"));
>+        stringBuilder.append("    threadgroup_barrier(mem_flags::mem_device);\n");
>+        stringBuilder.append("    threadgroup_barrier(mem_flags::mem_threadgroup);\n");
>+        stringBuilder.append("    threadgroup_barrier(mem_flags::mem_texture);\n");
>+        stringBuilder.append("}\n");
>+        return stringBuilder.toString();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "DeviceMemoryBarrierWithGroupSync") {
>+        ASSERT(!nativeFunctionDeclaration.parameters().size());
>+        stringBuilder.append(makeString("void ", outputFunctionName, "() {\n"));
>+        stringBuilder.append("    threadgroup_barrier(mem_flags::mem_device);\n");
>+        stringBuilder.append("}\n");
>+        return stringBuilder.toString();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "GroupMemoryBarrierWithGroupSync") {
>+        ASSERT(!nativeFunctionDeclaration.parameters().size());
>+        stringBuilder.append(makeString("void ", outputFunctionName, "() {\n"));
>+        stringBuilder.append("    threadgroup_barrier(mem_flags::mem_threadgroup);\n");
>+        stringBuilder.append("}\n");
>+        return stringBuilder.toString();
>+    }
>+
>+    if (nativeFunctionDeclaration.name().startsWith("Interlocked"_str)) {
>+        if (nativeFunctionDeclaration.name() == "InterlockedCompareExchange") {
>+            ASSERT(nativeFunctionDeclaration.parameters().size() == 4);
>+            ASSERT(is<AST::PointerType>(*nativeFunctionDeclaration.parameters()[0].type()));
>+            auto& firstArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[0].type());
>+            auto firstArgumentAddressSpace = firstArgumentPointer.addressSpace();
>+            auto firstArgumentPointee = getNativeName(firstArgumentPointer.elementType(), typeNamer);
>+            auto secondArgument = getNativeName(*nativeFunctionDeclaration.parameters()[1].type(), typeNamer);
>+            auto thirdArgument = getNativeName(*nativeFunctionDeclaration.parameters()[2].type(), typeNamer);
>+            ASSERT(is<AST::PointerType>(*nativeFunctionDeclaration.parameters()[3].type()));
>+            auto& fourthArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[3].type());
>+            auto fourthArgumentAddressSpace = fourthArgumentPointer.addressSpace();
>+            auto fourthArgumentPointee = getNativeName(fourthArgumentPointer.elementType(), typeNamer);
>+            stringBuilder.append(makeString("void ", outputFunctionName, '(', convertAddressSpace(firstArgumentAddressSpace), ' ', firstArgumentPointee, "* object, ", secondArgument, " compare, ", thirdArgument, " desired, ", convertAddressSpace(fourthArgumentAddressSpace), ' ', fourthArgumentPointee, "* out) {\n"));
>+            stringBuilder.append("    atomic_compare_exchange_weak_explicit(object, &compare, desired, memory_order_relaxed);\n");
>+            stringBuilder.append("    *out = compare;\n");
>+            stringBuilder.append("}\n");
>+            return stringBuilder.toString();
>+        }
>+
>+        ASSERT(nativeFunctionDeclaration.parameters().size() == 3);
>+        ASSERT(is<AST::PointerType>(*nativeFunctionDeclaration.parameters()[0].type()));
>+        auto& firstArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[0].type());
>+        auto firstArgumentAddressSpace = firstArgumentPointer.addressSpace();
>+        auto firstArgumentPointee = getNativeName(firstArgumentPointer.elementType(), typeNamer);
>+        auto secondArgument = getNativeName(*nativeFunctionDeclaration.parameters()[1].type(), typeNamer);
>+        ASSERT(is<AST::PointerType>(*nativeFunctionDeclaration.parameters()[2].type()));
>+        auto& thirdArgumentPointer = downcast<AST::PointerType>(*nativeFunctionDeclaration.parameters()[2].type());
>+        auto thirdArgumentAddressSpace = thirdArgumentPointer.addressSpace();
>+        auto thirdArgumentPointee = getNativeName(thirdArgumentPointer.elementType(), typeNamer);
>+        auto name = ([](String input) -> String {
>+            if (input == "Add")
>+                return "fetch_add"_str;
>+            if (input == "And")
>+                return "fetch_and"_str;
>+            if (input == "Exchange")
>+                return "exchange"_str;
>+            if (input == "Max")
>+                return "fetch_max"_str;
>+            if (input == "Min")
>+                return "fetch_min"_str;
>+            if (input == "Or")
>+                return "fetch_or"_str;
>+            ASSERT(input == "Xor");
>+            return "fetch_xor"_str;
>+        })(nativeFunctionDeclaration.name().substring("Interlocked"_str.length()));
>+        stringBuilder.append(makeString("void ", outputFunctionName, '(', convertAddressSpace(firstArgumentAddressSpace), ' ', firstArgumentPointee, "* object, ", secondArgument, " operand, ", convertAddressSpace(thirdArgumentAddressSpace), ' ', thirdArgumentPointee, "* out) {\n"));
>+        stringBuilder.append(makeString("    *out = atomic_fetch_", name, "_explicit(object, operand, memory_order_relaxed);\n"));
>+        stringBuilder.append("}\n");
>+        return stringBuilder.toString();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "Sample") {
>+        // FIXME: Implement this.
>+        CRASH();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "Load") {
>+        // FIXME: Implement this.
>+        CRASH();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "GetDimensions") {
>+        // FIXME: Implement this.
>+        CRASH();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "SampleBias") {
>+        // FIXME: Implement this.
>+        CRASH();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "SampleGrad") {
>+        // FIXME: Implement this.
>+        CRASH();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "SampleLevel") {
>+        // FIXME: Implement this.
>+        CRASH();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "Gather") {
>+        // FIXME: Implement this.
>+        CRASH();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "GatherRed") {
>+        // FIXME: Implement this.
>+        CRASH();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "SampleCmp") {
>+        // FIXME: Implement this.
>+        CRASH();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "SampleCmpLevelZero") {
>+        // FIXME: Implement this.
>+        CRASH();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "Store") {
>+        // FIXME: Implement this.
>+        CRASH();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "GatherAlpha") {
>+        // FIXME: Implement this.
>+        CRASH();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "GatherBlue") {
>+        // FIXME: Implement this.
>+        CRASH();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "GatherCmp") {
>+        // FIXME: Implement this.
>+        CRASH();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "GatherCmpRed") {
>+        // FIXME: Implement this.
>+        CRASH();
>+    }
>+
>+    if (nativeFunctionDeclaration.name() == "GatherGreen") {
>+        // FIXME: Implement this.
>+        CRASH();
>+    }
>+
>+    // FIXME: Add all the functions that the compiler generated.
>+
>+    ASSERT_NOT_REACHED();
>+    return String();
>+}
>+
>+} // namespace Metal
>+
>+} // namespace WHLSL
>+
>+} // namespace WebCore
>+
>+#endif
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLNativeFunctionWriter.h b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLNativeFunctionWriter.h
>new file mode 100644
>index 00000000000..2e9bd4e3cbd
>--- /dev/null
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLNativeFunctionWriter.h
>@@ -0,0 +1,54 @@
>+/*
>+ * Copyright (C) 2019 Apple Inc. All rights reserved.
>+ *
>+ * Redistribution and use in source and binary forms, with or without
>+ * modification, are permitted provided that the following conditions
>+ * are met:
>+ * 1. Redistributions of source code must retain the above copyright
>+ *    notice, this list of conditions and the following disclaimer.
>+ * 2. Redistributions in binary form must reproduce the above copyright
>+ *    notice, this list of conditions and the following disclaimer in the
>+ *    documentation and/or other materials provided with the distribution.
>+ *
>+ * THIS SOFTWARE IS PROVIDED BY APPLE INC. AND ITS CONTRIBUTORS ``AS IS''
>+ * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO,
>+ * THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
>+ * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL APPLE INC. OR ITS CONTRIBUTORS
>+ * BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
>+ * CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
>+ * SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
>+ * INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
>+ * CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
>+ * ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
>+ * THE POSSIBILITY OF SUCH DAMAGE.
>+ */
>+
>+#pragma once
>+
>+#if ENABLE(WEBGPU)
>+
>+#include <wtf/text/WTFString.h>
>+
>+namespace WebCore {
>+
>+namespace WHLSL {
>+
>+namespace AST {
>+
>+class NativeFunctionDeclaration;
>+
>+}
>+
>+namespace Metal {
>+
>+class TypeNamer;
>+
>+String writeNativeFunction(AST::NativeFunctionDeclaration&, String& outputFunctionName, TypeNamer&);
>+
>+}
>+
>+}
>+
>+}
>+
>+#endif
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLTypeNamer.cpp b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLTypeNamer.cpp
>new file mode 100644
>index 00000000000..2318eb8676a
>--- /dev/null
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLTypeNamer.cpp
>@@ -0,0 +1,706 @@
>+/*
>+ * Copyright (C) 2019 Apple Inc. All rights reserved.
>+ *
>+ * Redistribution and use in source and binary forms, with or without
>+ * modification, are permitted provided that the following conditions
>+ * are met:
>+ * 1. Redistributions of source code must retain the above copyright
>+ *    notice, this list of conditions and the following disclaimer.
>+ * 2. Redistributions in binary form must reproduce the above copyright
>+ *    notice, this list of conditions and the following disclaimer in the
>+ *    documentation and/or other materials provided with the distribution.
>+ *
>+ * THIS SOFTWARE IS PROVIDED BY APPLE INC. AND ITS CONTRIBUTORS ``AS IS''
>+ * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO,
>+ * THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
>+ * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL APPLE INC. OR ITS CONTRIBUTORS
>+ * BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
>+ * CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
>+ * SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
>+ * INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
>+ * CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
>+ * ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
>+ * THE POSSIBILITY OF SUCH DAMAGE.
>+ */
>+
>+#include "config.h"
>+#include "WHLSLTypeNamer.h"
>+
>+#if ENABLE(WEBGPU)
>+
>+#include "WHLSLAddressSpace.h"
>+#include "WHLSLArrayReferenceType.h"
>+#include "WHLSLArrayType.h"
>+#include "WHLSLEnumerationDefinition.h"
>+#include "WHLSLEnumerationMember.h"
>+#include "WHLSLNativeTypeDeclaration.h"
>+#include "WHLSLPointerType.h"
>+#include "WHLSLStructureDefinition.h"
>+#include "WHLSLTypeDefinition.h"
>+#include "WHLSLTypeReference.h"
>+#include "WHLSLVisitor.h"
>+#include <algorithm>
>+#include <functional>
>+#include <wtf/HashMap.h>
>+#include <wtf/HashSet.h>
>+#include <wtf/Optional.h>
>+#include <wtf/UniqueRef.h>
>+#include <wtf/Vector.h>
>+#include <wtf/text/StringBuilder.h>
>+#include <wtf/text/StringConcatenateNumbers.h>
>+
>+namespace WebCore {
>+
>+namespace WHLSL {
>+
>+namespace Metal {
>+
>+class BaseTypeNameNode {
>+public:
>+    BaseTypeNameNode(BaseTypeNameNode* parent, String&& mangledName)
>+        : m_parent(parent)
>+        , m_mangledName(mangledName)
>+    {
>+    }
>+    virtual ~BaseTypeNameNode() = default;
>+    virtual bool isArrayTypeNameNode() const { return false; }
>+    virtual bool isArrayReferenceTypeNameNode() const { return false; }
>+    virtual bool isPointerTypeNameNode() const { return false; }
>+    virtual bool isReferenceTypeNameNode() const { return false; }
>+    Vector<UniqueRef<BaseTypeNameNode>>& children() { return m_children; }
>+    void append(UniqueRef<BaseTypeNameNode>&& child)
>+    {
>+        m_children.append(WTFMove(child));
>+    }
>+    BaseTypeNameNode* parent() { return m_parent; }
>+    const String& mangledName() const { return m_mangledName; }
>+
>+private:
>+    Vector<UniqueRef<BaseTypeNameNode>> m_children;
>+    BaseTypeNameNode* m_parent;
>+    String m_mangledName;
>+};
>+
>+class ArrayTypeNameNode : public BaseTypeNameNode {
>+public:
>+    ArrayTypeNameNode(BaseTypeNameNode* parent, String&& mangledName, unsigned numElements)
>+        : BaseTypeNameNode(parent, WTFMove(mangledName))
>+        , m_numElements(numElements)
>+    {
>+    }
>+    virtual ~ArrayTypeNameNode() = default;
>+    bool isArrayTypeNameNode() const override { return true; }
>+    unsigned numElements() const { return m_numElements; }
>+
>+private:
>+    unsigned m_numElements;
>+};
>+
>+class ArrayReferenceTypeNameNode : public BaseTypeNameNode {
>+public:
>+    ArrayReferenceTypeNameNode(BaseTypeNameNode* parent, String&& mangledName, AST::AddressSpace addressSpace)
>+        : BaseTypeNameNode(parent, WTFMove(mangledName))
>+        , m_addressSpace(addressSpace)
>+    {
>+    }
>+    virtual ~ArrayReferenceTypeNameNode() = default;
>+    bool isArrayReferenceTypeNameNode() const override { return true; }
>+    AST::AddressSpace addressSpace() const { return m_addressSpace; }
>+
>+private:
>+    AST::AddressSpace m_addressSpace;
>+};
>+
>+class PointerTypeNameNode : public BaseTypeNameNode {
>+public:
>+    PointerTypeNameNode(BaseTypeNameNode* parent, String&& mangledName, AST::AddressSpace addressSpace)
>+        : BaseTypeNameNode(parent, WTFMove(mangledName))
>+        , m_addressSpace(addressSpace)
>+    {
>+    }
>+    virtual ~PointerTypeNameNode() = default;
>+    bool isPointerTypeNameNode() const override { return true; }
>+    AST::AddressSpace addressSpace() const { return m_addressSpace; }
>+
>+private:
>+    AST::AddressSpace m_addressSpace;
>+};
>+
>+class ReferenceTypeNameNode : public BaseTypeNameNode {
>+public:
>+    ReferenceTypeNameNode(BaseTypeNameNode* parent, String&& mangledName, AST::NamedType& namedType)
>+        : BaseTypeNameNode(parent, WTFMove(mangledName))
>+        , m_namedType(namedType)
>+    {
>+    }
>+    virtual ~ReferenceTypeNameNode() = default;
>+    bool isReferenceTypeNameNode() const override { return true; }
>+    AST::NamedType& namedType() { return m_namedType; }
>+
>+private:
>+    AST::NamedType& m_namedType;
>+};
>+
>+}
>+
>+}
>+
>+}
>+
>+#define SPECIALIZE_TYPE_TRAITS_WHLSL_BASE_TYPE_NAMED_NODE(ToValueTypeName, predicate) \
>+SPECIALIZE_TYPE_TRAITS_BEGIN(WebCore::WHLSL::Metal::ToValueTypeName) \
>+    static bool isType(const WebCore::WHLSL::Metal::BaseTypeNameNode& type) { return type.predicate; } \
>+SPECIALIZE_TYPE_TRAITS_END()
>+
>+SPECIALIZE_TYPE_TRAITS_WHLSL_BASE_TYPE_NAMED_NODE(ArrayTypeNameNode, isArrayTypeNameNode())
>+
>+SPECIALIZE_TYPE_TRAITS_WHLSL_BASE_TYPE_NAMED_NODE(ArrayReferenceTypeNameNode, isArrayReferenceTypeNameNode())
>+
>+SPECIALIZE_TYPE_TRAITS_WHLSL_BASE_TYPE_NAMED_NODE(PointerTypeNameNode, isPointerTypeNameNode())
>+
>+SPECIALIZE_TYPE_TRAITS_WHLSL_BASE_TYPE_NAMED_NODE(ReferenceTypeNameNode, isReferenceTypeNameNode())
>+
>+namespace WebCore {
>+
>+namespace WHLSL {
>+
>+namespace Metal {
>+
>+TypeNamer::TypeNamer(Program& program)
>+    : m_program(program)
>+{
>+}
>+
>+TypeNamer::~TypeNamer() = default;
>+
>+void TypeNamer::visit(AST::UnnamedType& unnamedType)
>+{
>+    insert(unnamedType, m_trie);
>+}
>+
>+void TypeNamer::visit(AST::EnumerationDefinition& enumerationDefinition)
>+{
>+    auto addResult = m_namedTypeMapping.add(&enumerationDefinition, generateNextTypeName());
>+    ASSERT_UNUSED(addResult, addResult.isNewEntry);
>+    for (auto& enumerationMember : enumerationDefinition.enumerationMembers()) {
>+        auto addResult = m_enumerationMemberMapping.add(&static_cast<AST::EnumerationMember&>(enumerationMember), generateNextEnumerationMemberName());
>+        ASSERT_UNUSED(addResult, addResult.isNewEntry);
>+    }
>+    Visitor::visit(enumerationDefinition);
>+}
>+
>+void TypeNamer::visit(AST::NativeTypeDeclaration&)
>+{
>+    // Native type declarations already have names, and are already declared in Metal.
>+}
>+
>+static Vector<UniqueRef<BaseTypeNameNode>>::iterator findInVector(AST::UnnamedType& unnamedType, Vector<UniqueRef<BaseTypeNameNode>>& types)
>+{
>+    return std::find_if(types.begin(), types.end(), [&](BaseTypeNameNode& baseTypeNameNode) -> bool {
>+        if (is<AST::TypeReference>(unnamedType) && is<ReferenceTypeNameNode>(baseTypeNameNode)) {
>+            auto* resolvedType = downcast<AST::TypeReference>(unnamedType).resolvedType();
>+            ASSERT(resolvedType);
>+            return resolvedType == &downcast<ReferenceTypeNameNode>(baseTypeNameNode).namedType();
>+        }
>+        if (is<AST::PointerType>(unnamedType) && is<PointerTypeNameNode>(baseTypeNameNode))
>+            return downcast<AST::PointerType>(unnamedType).addressSpace() == downcast<PointerTypeNameNode>(baseTypeNameNode).addressSpace();
>+        if (is<AST::ArrayReferenceType>(unnamedType) && is<ArrayReferenceTypeNameNode>(baseTypeNameNode))
>+            return downcast<AST::ArrayReferenceType>(unnamedType).addressSpace() == downcast<ArrayReferenceTypeNameNode>(baseTypeNameNode).addressSpace();
>+        if (is<AST::ArrayType>(unnamedType) && is<ArrayTypeNameNode>(baseTypeNameNode))
>+            return downcast<AST::ArrayType>(unnamedType).numElements() == downcast<ArrayTypeNameNode>(baseTypeNameNode).numElements();
>+        return false;
>+    });
>+}
>+
>+static BaseTypeNameNode& find(AST::UnnamedType& unnamedType, Vector<UniqueRef<BaseTypeNameNode>>& types)
>+{
>+    auto& vectorToSearch = ([&]() -> Vector<UniqueRef<BaseTypeNameNode>>& {
>+        if (is<AST::TypeReference>(unnamedType))
>+            return types;
>+        if (is<AST::PointerType>(unnamedType))
>+            return find(downcast<AST::PointerType>(unnamedType).elementType(), types).children();
>+        if (is<AST::ArrayReferenceType>(unnamedType))
>+            return find(downcast<AST::ArrayReferenceType>(unnamedType).elementType(), types).children();
>+        ASSERT(is<AST::ArrayType>(unnamedType));
>+        return find(downcast<AST::ArrayType>(unnamedType).type(), types).children();
>+    })();
>+    auto iterator = findInVector(unnamedType, vectorToSearch);
>+    ASSERT(iterator != types.end());
>+    return *iterator;
>+}
>+
>+void TypeNamer::visit(AST::StructureDefinition& structureDefinition)
>+{
>+    {
>+        auto addResult = m_namedTypeMapping.add(&structureDefinition, generateNextTypeName());
>+        ASSERT_UNUSED(addResult, addResult.isNewEntry);
>+    }
>+    Visitor::visit(structureDefinition);
>+    {
>+        Vector<std::reference_wrapper<BaseTypeNameNode>> neighbors;
>+        for (auto& structureElement : structureDefinition.structureElements()) {
>+            auto addResult = m_structureElementMapping.add(&structureElement, generateNextStructureElementName());
>+            ASSERT_UNUSED(addResult, addResult.isNewEntry);
>+            neighbors.append(find(structureElement.type(), m_trie));
>+        }
>+        auto addResult = m_dependencyGraph.add(&structureDefinition, WTFMove(neighbors));
>+        ASSERT_UNUSED(addResult, addResult.isNewEntry);
>+    }
>+}
>+
>+void TypeNamer::visit(AST::TypeDefinition& typeDefinition)
>+{
>+    {
>+        auto addResult = m_namedTypeMapping.add(&typeDefinition, generateNextTypeName());
>+        ASSERT_UNUSED(addResult, addResult.isNewEntry);
>+    }
>+    Visitor::visit(typeDefinition);
>+    {
>+        Vector<std::reference_wrapper<BaseTypeNameNode>> neighbors = { find(typeDefinition.type(), m_trie) };
>+        auto addResult = m_dependencyGraph.add(&typeDefinition, WTFMove(neighbors));
>+        ASSERT_UNUSED(addResult, addResult.isNewEntry);
>+    }
>+}
>+
>+String TypeNamer::mangledNameForType(AST::NativeTypeDeclaration& nativeTypeDeclaration)
>+{
>+    if (nativeTypeDeclaration.name() == "void")
>+        return "void"_str;
>+    if (nativeTypeDeclaration.name() == "bool")
>+        return "bool"_str;
>+    if (nativeTypeDeclaration.name() == "uchar")
>+        return "uint8_t"_str;
>+    if (nativeTypeDeclaration.name() == "ushort")
>+        return "uint16_t"_str;
>+    if (nativeTypeDeclaration.name() == "uint")
>+        return "uint32_t"_str;
>+    if (nativeTypeDeclaration.name() == "char")
>+        return "int8_t"_str;
>+    if (nativeTypeDeclaration.name() == "short")
>+        return "int16_t"_str;
>+    if (nativeTypeDeclaration.name() == "int")
>+        return "int32_t"_str;
>+    if (nativeTypeDeclaration.name() == "half")
>+        return "half"_str;
>+    if (nativeTypeDeclaration.name() == "float")
>+        return "float"_str;
>+    if (nativeTypeDeclaration.name() == "atomic_int")
>+        return "atomic_int"_str;
>+    if (nativeTypeDeclaration.name() == "atomic_uint")
>+        return "atomic_uint"_str;
>+    if (nativeTypeDeclaration.name() == "sampler")
>+        return "sampler"_str;
>+    if (nativeTypeDeclaration.name() == "vector") {
>+        ASSERT(nativeTypeDeclaration.typeArguments().size() == 2);
>+        ASSERT(WTF::holds_alternative<UniqueRef<AST::TypeReference>>(nativeTypeDeclaration.typeArguments()[0]));
>+        auto& typeReference = WTF::get<UniqueRef<AST::TypeReference>>(nativeTypeDeclaration.typeArguments()[0]);
>+        auto& unifyNode = typeReference->unifyNode();
>+        ASSERT(is<AST::NamedType>(unifyNode));
>+        auto& namedType = downcast<AST::NamedType>(unifyNode);
>+        ASSERT(is<AST::NativeTypeDeclaration>(namedType));
>+        auto& parameterType = downcast<AST::NativeTypeDeclaration>(namedType);
>+        auto prefix = ([&]() -> String {
>+            if (parameterType.name() == "bool")
>+                return "bool";
>+            if (parameterType.name() == "uchar")
>+                return "uchar";
>+            if (parameterType.name() == "ushort")
>+                return "ushort";
>+            if (parameterType.name() == "uint")
>+                return "uint";
>+            if (parameterType.name() == "char")
>+                return "char";
>+            if (parameterType.name() == "short")
>+                return "short";
>+            if (parameterType.name() == "int")
>+                return "int";
>+            if (parameterType.name() == "half")
>+                return "half";
>+            ASSERT(parameterType.name() == "float");
>+            return "float";
>+        })();
>+        ASSERT(WTF::holds_alternative<AST::ConstantExpression>(nativeTypeDeclaration.typeArguments()[1]));
>+        auto& constantExpression = WTF::get<AST::ConstantExpression>(nativeTypeDeclaration.typeArguments()[0]);
>+        auto& integerLiteral = constantExpression.integerLiteral();
>+        auto suffix = ([&]() -> String {
>+            switch (integerLiteral.value()) {
>+            case 2:
>+                return "2"_str;
>+            case 3:
>+                return "3"_str;
>+            default:
>+                ASSERT(integerLiteral.value() == 4);
>+                return "4"_str;
>+            }
>+        })();
>+        return makeString(prefix, suffix);
>+    }
>+    if (nativeTypeDeclaration.name() == "matrix") {
>+        ASSERT(nativeTypeDeclaration.typeArguments().size() == 3);
>+        ASSERT(WTF::holds_alternative<UniqueRef<AST::TypeReference>>(nativeTypeDeclaration.typeArguments()[0]));
>+        auto& typeReference = WTF::get<UniqueRef<AST::TypeReference>>(nativeTypeDeclaration.typeArguments()[0]);
>+        auto& unifyNode = typeReference->unifyNode();
>+        ASSERT(is<AST::NamedType>(unifyNode));
>+        auto& namedType = downcast<AST::NamedType>(unifyNode);
>+        ASSERT(is<AST::NativeTypeDeclaration>(namedType));
>+        auto& parameterType = downcast<AST::NativeTypeDeclaration>(namedType);
>+        auto prefix = ([&]() -> String {
>+            if (parameterType.name() == "half")
>+                return "half";
>+            ASSERT(parameterType.name() == "float");
>+            return "float";
>+        })();
>+        ASSERT(WTF::holds_alternative<AST::ConstantExpression>(nativeTypeDeclaration.typeArguments()[1]));
>+        auto& constantExpression1 = WTF::get<AST::ConstantExpression>(nativeTypeDeclaration.typeArguments()[0]);
>+        auto& integerLiteral1 = constantExpression1.integerLiteral();
>+        auto middle = ([&]() -> String {
>+            switch (integerLiteral1.value()) {
>+            case 2:
>+                return "2"_str;
>+            case 3:
>+                return "3"_str;
>+            default:
>+                ASSERT(integerLiteral1.value() == 4);
>+                return "4"_str;
>+            }
>+        })();
>+        ASSERT(WTF::holds_alternative<AST::ConstantExpression>(nativeTypeDeclaration.typeArguments()[2]));
>+        auto& constantExpression2 = WTF::get<AST::ConstantExpression>(nativeTypeDeclaration.typeArguments()[0]);
>+        auto& integerLiteral2 = constantExpression2.integerLiteral();
>+        auto suffix = ([&]() -> String {
>+            switch (integerLiteral2.value()) {
>+            case 2:
>+                return "2"_str;
>+            case 3:
>+                return "3"_str;
>+            default:
>+                ASSERT(integerLiteral2.value() == 4);
>+                return "4"_str;
>+            }
>+        })();
>+        return makeString(prefix, middle, 'x', suffix);
>+    }
>+    ASSERT(nativeTypeDeclaration.typeArguments().size() == 1);
>+    ASSERT(WTF::holds_alternative<UniqueRef<AST::TypeReference>>(nativeTypeDeclaration.typeArguments()[0]));
>+    auto& typeReference = WTF::get<UniqueRef<AST::TypeReference>>(nativeTypeDeclaration.typeArguments()[0]);
>+    auto prefix = ([&]() -> String {
>+        if (nativeTypeDeclaration.name() == "Texture1D")
>+            return "texture1d"_str;
>+        if (nativeTypeDeclaration.name() == "RWTexture1D")
>+            return "texture1d"_str;
>+        if (nativeTypeDeclaration.name() == "Texture1DArray")
>+            return "texture1d_array"_str;
>+        if (nativeTypeDeclaration.name() == "RWTexture1DArray")
>+            return "texture1d_array"_str;
>+        if (nativeTypeDeclaration.name() == "Texture2D")
>+            return "texture2d"_str;
>+        if (nativeTypeDeclaration.name() == "RWTexture2D")
>+            return "texture2d"_str;
>+        if (nativeTypeDeclaration.name() == "Texture2DArray")
>+            return "texture2d_array"_str;
>+        if (nativeTypeDeclaration.name() == "RWTexture2DArray")
>+            return "texture2d_array"_str;
>+        if (nativeTypeDeclaration.name() == "Texture3D")
>+            return "texture3d"_str;
>+        if (nativeTypeDeclaration.name() == "RWTexture3D")
>+            return "texture3d"_str;
>+        if (nativeTypeDeclaration.name() == "TextureCube")
>+            return "texturecube"_str;
>+        if (nativeTypeDeclaration.name() == "TextureDepth2D")
>+            return "depth2d"_str;
>+        if (nativeTypeDeclaration.name() == "RWTextureDepth2D")
>+            return "depth2d"_str;
>+        if (nativeTypeDeclaration.name() == "TextureDepth2DArray")
>+            return "depth2d_array"_str;
>+        if (nativeTypeDeclaration.name() == "RWTextureDepth2DArray")
>+            return "depth2d_array"_str;
>+        ASSERT(nativeTypeDeclaration.name() == "TextureDepthCube");
>+        return "depthcube"_str;
>+    })();
>+    auto innerType = ([&]() -> String {
>+        if (typeReference->name() == "ushort")
>+            return "ushort"_str;
>+        if (typeReference->name() == "ushort2")
>+            return "ushort"_str;
>+        if (typeReference->name() == "ushort3")
>+            return "ushort"_str;
>+        if (typeReference->name() == "ushort4")
>+            return "ushort"_str;
>+        if (typeReference->name() == "uint")
>+            return "uint"_str;
>+        if (typeReference->name() == "uint2")
>+            return "uint"_str;
>+        if (typeReference->name() == "uint3")
>+            return "uint"_str;
>+        if (typeReference->name() == "uint4")
>+            return "uint"_str;
>+        if (typeReference->name() == "short")
>+            return "short"_str;
>+        if (typeReference->name() == "short2")
>+            return "short"_str;
>+        if (typeReference->name() == "short3")
>+            return "short"_str;
>+        if (typeReference->name() == "short4")
>+            return "short"_str;
>+        if (typeReference->name() == "int")
>+            return "int"_str;
>+        if (typeReference->name() == "int2")
>+            return "int"_str;
>+        if (typeReference->name() == "int3")
>+            return "int"_str;
>+        if (typeReference->name() == "int4")
>+            return "int"_str;
>+        if (typeReference->name() == "half")
>+            return "half"_str;
>+        if (typeReference->name() == "half2")
>+            return "half"_str;
>+        if (typeReference->name() == "half3")
>+            return "half"_str;
>+        if (typeReference->name() == "half4")
>+            return "half"_str;
>+        if (typeReference->name() == "float")
>+            return "float"_str;
>+        if (typeReference->name() == "float2")
>+            return "float"_str;
>+        if (typeReference->name() == "float3")
>+            return "float"_str;
>+        ASSERT(typeReference->name() == "float4");
>+        return "float"_str;
>+    })();
>+    // FIXME: Specify the second template argument to Metal texture types.
>+    return makeString(prefix, '<', innerType, '>');
>+}
>+
>+UniqueRef<BaseTypeNameNode> TypeNamer::createNameNode(AST::UnnamedType& unnamedType, BaseTypeNameNode* parent)
>+{
>+    if (is<AST::TypeReference>(unnamedType)) {
>+        auto& typeReference = downcast<AST::TypeReference>(unnamedType);
>+        ASSERT(typeReference.resolvedType());
>+        return makeUniqueRef<ReferenceTypeNameNode>(parent, generateNextTypeName(), *typeReference.resolvedType());
>+    }
>+    if (is<AST::PointerType>(unnamedType)) {
>+        auto& pointerType = downcast<AST::PointerType>(unnamedType);
>+        return makeUniqueRef<PointerTypeNameNode>(parent, generateNextTypeName(), pointerType.addressSpace());
>+    }
>+    if (is<AST::ArrayReferenceType>(unnamedType)) {
>+        auto& arrayReferenceType = downcast<AST::ArrayReferenceType>(unnamedType);
>+        return makeUniqueRef<ArrayReferenceTypeNameNode>(parent, generateNextTypeName(), arrayReferenceType.addressSpace());
>+    }
>+    ASSERT(is<AST::ArrayType>(unnamedType));
>+    auto& arrayType = downcast<AST::ArrayType>(unnamedType);
>+    return makeUniqueRef<ArrayTypeNameNode>(parent, generateNextTypeName(), arrayType.numElements());
>+}
>+
>+size_t TypeNamer::insert(AST::UnnamedType& unnamedType, Vector<UniqueRef<BaseTypeNameNode>>& types)
>+{
>+    Vector<UniqueRef<BaseTypeNameNode>>* vectorToInsertInto { nullptr };
>+    BaseTypeNameNode* parent { nullptr };
>+    if (is<AST::TypeReference>(unnamedType)) {
>+        vectorToInsertInto = &types;
>+        parent = nullptr;
>+    } else if (is<AST::PointerType>(unnamedType)) {
>+        auto& item = types[insert(downcast<AST::PointerType>(unnamedType).elementType(), types)];
>+        vectorToInsertInto = &item->children();
>+        parent = &item;
>+    } else if (is<AST::ArrayReferenceType>(unnamedType)) {
>+        auto& item = types[insert(downcast<AST::ArrayReferenceType>(unnamedType).elementType(), types)];
>+        vectorToInsertInto = &item->children();
>+        parent = &item;
>+    } else {
>+        ASSERT(is<AST::ArrayType>(unnamedType));
>+        auto& item = types[insert(downcast<AST::ArrayType>(unnamedType).type(), types)];
>+        vectorToInsertInto = &item->children();
>+        parent = &item;
>+    }
>+    ASSERT(vectorToInsertInto);
>+
>+    auto iterator = findInVector(unnamedType, *vectorToInsertInto);
>+    if (iterator == vectorToInsertInto->end()) {
>+        auto result = createNameNode(unnamedType, parent);
>+        {
>+            auto addResult = m_unnamedTypeMapping.add(&unnamedType, &result);
>+            ASSERT_UNUSED(addResult, addResult.isNewEntry);
>+        }
>+        vectorToInsertInto->append(WTFMove(result));
>+        return vectorToInsertInto->size() - 1;
>+    }
>+    auto addResult = m_unnamedTypeMapping.add(&unnamedType, &*iterator);
>+    ASSERT_UNUSED(addResult, addResult.isNewEntry);
>+    return iterator - vectorToInsertInto->begin();
>+}
>+
>+class MetalTypeDeclarationWriter : public Visitor {
>+public:
>+    MetalTypeDeclarationWriter(std::function<String(AST::NamedType&)>&& mangledNameForNamedType, std::function<String(AST::UnnamedType&)>&& mangledNameForUnnamedType, std::function<String(AST::EnumerationMember&)>&& mangledNameForEnumerationMember)
>+        : m_mangledNameForNamedType(WTFMove(mangledNameForNamedType))
>+        , m_mangledNameForUnnamedType(WTFMove(mangledNameForUnnamedType))
>+        , m_mangledNameForEnumerationMember(WTFMove(mangledNameForEnumerationMember))
>+    {
>+    }
>+
>+    String toString() { return m_stringBuilder.toString(); }
>+
>+private:
>+    void visit(AST::EnumerationDefinition& enumerationDefinition)
>+    {
>+        auto& baseType = enumerationDefinition.type().unifyNode();
>+        ASSERT(is<AST::NamedType>(baseType));
>+        m_stringBuilder.append(makeString("enum class ", m_mangledNameForNamedType(enumerationDefinition), " : ", m_mangledNameForNamedType(downcast<AST::NamedType>(baseType)), " {\n"));
>+        for (auto& enumerationMember : enumerationDefinition.enumerationMembers())
>+            m_stringBuilder.append(makeString("    ", m_mangledNameForEnumerationMember(enumerationMember), ",\n"));
>+        m_stringBuilder.append("};\n");
>+    }
>+
>+    void visit(AST::StructureDefinition& structureDefinition)
>+    {
>+        m_stringBuilder.append(makeString("struct ", m_mangledNameForNamedType(structureDefinition), ";\n"));
>+    }
>+
>+    std::function<String(AST::NamedType&)> m_mangledNameForNamedType;
>+    std::function<String(AST::UnnamedType&)> m_mangledNameForUnnamedType;
>+    std::function<String(AST::EnumerationMember&)>&& m_mangledNameForEnumerationMember;
>+    StringBuilder m_stringBuilder;
>+};
>+
>+String TypeNamer::metalTypeDeclarations()
>+{
>+    MetalTypeDeclarationWriter metalTypeDeclarationWriter([&](AST::NamedType& namedType) -> String {
>+        return mangledNameForType(namedType);
>+    }, [&](AST::UnnamedType& unnamedType) -> String {
>+        return mangledNameForType(unnamedType);
>+    }, [&](AST::EnumerationMember& enumerationMember) -> String {
>+        return mangledNameForEnumerationMember(enumerationMember);
>+    });
>+    metalTypeDeclarationWriter.Visitor::visit(m_program);
>+    return metalTypeDeclarationWriter.toString();
>+}
>+
>+static String toString(AST::AddressSpace addressSpace)
>+{
>+    switch (addressSpace) {
>+    case AST::AddressSpace::Constant:
>+        return "constant"_str;
>+    case AST::AddressSpace::Device:
>+        return "device"_str;
>+    case AST::AddressSpace::Threadgroup:
>+        return "threadgroup"_str;
>+    default:
>+        ASSERT(addressSpace == AST::AddressSpace::Thread);
>+        return "thread"_str;
>+    }
>+}
>+
>+void TypeNamer::emitUnnamedTypeDefinition(BaseTypeNameNode& baseTypeNameNode, HashSet<AST::NamedType*>& emittedNamedTypes, HashSet<BaseTypeNameNode*>& emittedUnnamedTypes, StringBuilder& stringBuilder)
>+{
>+    if (emittedUnnamedTypes.contains(&baseTypeNameNode))
>+        return;
>+    if (baseTypeNameNode.parent())
>+        emitUnnamedTypeDefinition(*baseTypeNameNode.parent(), emittedNamedTypes, emittedUnnamedTypes, stringBuilder);
>+    if (is<ReferenceTypeNameNode>(baseTypeNameNode)) {
>+        auto& namedType = downcast<ReferenceTypeNameNode>(baseTypeNameNode).namedType();
>+        emitNamedTypeDefinition(namedType, emittedNamedTypes, emittedUnnamedTypes, stringBuilder);
>+        stringBuilder.append(makeString("typedef ", mangledNameForType(namedType), ' ', baseTypeNameNode.mangledName(), ";\n"));
>+    } else if (is<PointerTypeNameNode>(baseTypeNameNode)) {
>+        auto& pointerType = downcast<PointerTypeNameNode>(baseTypeNameNode);
>+        ASSERT(baseTypeNameNode.parent());
>+        stringBuilder.append(makeString("typedef ", toString(pointerType.addressSpace()), " ", pointerType.parent()->mangledName(), "* ", pointerType.mangledName(), ";\n"));
>+    } else if (is<ArrayReferenceTypeNameNode>(baseTypeNameNode)) {
>+        auto& arrayReferenceType = downcast<ArrayReferenceTypeNameNode>(baseTypeNameNode);
>+        ASSERT(baseTypeNameNode.parent());
>+        stringBuilder.append(makeString("struct ", arrayReferenceType.mangledName(), "{ \n"));
>+        stringBuilder.append(makeString("    ", toString(arrayReferenceType.addressSpace()), " ", arrayReferenceType.parent()->mangledName(), "* pointer;\n"));
>+        stringBuilder.append("    unsigned length;\n");
>+        stringBuilder.append("};\n");
>+    } else {
>+        ASSERT(is<ArrayTypeNameNode>(baseTypeNameNode));
>+        auto& arrayType = downcast<ArrayTypeNameNode>(baseTypeNameNode);
>+        ASSERT(baseTypeNameNode.parent());
>+        stringBuilder.append(makeString("typedef Array<", arrayType.parent()->mangledName(), ", ", arrayType.numElements(), "> ", arrayType.mangledName(), ";\n"));
>+    }
>+    emittedUnnamedTypes.add(&baseTypeNameNode);
>+}
>+
>+void TypeNamer::emitNamedTypeDefinition(AST::NamedType& namedType, HashSet<AST::NamedType*>& emittedNamedTypes, HashSet<BaseTypeNameNode*>& emittedUnnamedTypes, StringBuilder& stringBuilder)
>+{
>+    if (emittedNamedTypes.contains(&namedType))
>+        return;
>+    auto iterator = m_dependencyGraph.find(&namedType);
>+    ASSERT(iterator != m_dependencyGraph.end());
>+    for (auto& baseTypeNameNode : iterator->value)
>+        emitUnnamedTypeDefinition(baseTypeNameNode, emittedNamedTypes, emittedUnnamedTypes, stringBuilder);
>+    if (is<AST::EnumerationDefinition>(namedType)) {
>+        // We already emitted this in the type declaration section. There's nothing to do.
>+    } else if (is<AST::NativeTypeDeclaration>(namedType)) {
>+        // Native types already have definitions. There's nothing to do.
>+    } else if (is<AST::StructureDefinition>(namedType)) {
>+        auto& structureDefinition = downcast<AST::StructureDefinition>(namedType);
>+        stringBuilder.append(makeString("struct ", mangledNameForType(structureDefinition), " {\n"));
>+        for (auto& structureElement : structureDefinition.structureElements())
>+            stringBuilder.append(makeString("    ", mangledNameForType(structureElement.type()), ' ', mangledNameForStructureElement(structureElement), ";\n"));
>+        stringBuilder.append("};\n");
>+    } else {
>+        ASSERT(is<AST::TypeDefinition>(namedType));
>+        auto& typeDefinition = downcast<AST::TypeDefinition>(namedType);
>+        stringBuilder.append(makeString("typedef ", mangledNameForType(typeDefinition.type()), ' ', mangledNameForType(typeDefinition), ";\n"));
>+    }
>+    emittedNamedTypes.add(&namedType);
>+}
>+
>+String TypeNamer::metalTypeDefinitions()
>+{
>+    HashSet<AST::NamedType*> emittedNamedTypes;
>+    HashSet<BaseTypeNameNode*> emittedUnnamedTypes;
>+    StringBuilder stringBuilder;
>+    for (auto& keyValuePair : m_dependencyGraph)
>+        emitNamedTypeDefinition(*keyValuePair.key, emittedNamedTypes, emittedUnnamedTypes, stringBuilder);
>+    for (auto& baseTypeNameNode : m_trie)
>+        emitUnnamedTypeDefinition(baseTypeNameNode, emittedNamedTypes, emittedUnnamedTypes, stringBuilder);
>+    return stringBuilder.toString();
>+}
>+
>+String TypeNamer::mangledNameForType(AST::UnnamedType& unnamedType)
>+{
>+    return find(unnamedType, m_trie).mangledName();
>+}
>+
>+String TypeNamer::mangledNameForType(AST::NamedType& namedType)
>+{
>+    if (is<AST::NativeTypeDeclaration>(namedType))
>+        return mangledNameForType(downcast<AST::NativeTypeDeclaration>(namedType));
>+    auto iterator = m_namedTypeMapping.find(&namedType);
>+    ASSERT(iterator != m_namedTypeMapping.end());
>+    return iterator->value;
>+}
>+
>+
>+String TypeNamer::mangledNameForEnumerationMember(AST::EnumerationMember& enumerationMember)
>+{
>+    auto iterator = m_enumerationMemberMapping.find(&enumerationMember);
>+    ASSERT(iterator != m_enumerationMemberMapping.end());
>+    return iterator->value;
>+}
>+
>+String TypeNamer::mangledNameForStructureElement(AST::StructureElement& structureElement)
>+{
>+    auto iterator = m_structureElementMapping.find(&structureElement);
>+    ASSERT(iterator != m_structureElementMapping.end());
>+    return iterator->value;
>+}
>+
>+String TypeNamer::metalTypes()
>+{
>+    Visitor::visit(m_program);
>+    StringBuilder stringBuilder;
>+    stringBuilder.append(metalTypeDeclarations());
>+    stringBuilder.append('\n');
>+    stringBuilder.append(metalTypeDefinitions());
>+    return stringBuilder.toString();
>+}
>+
>+} // namespace Metal
>+
>+} // namespace WHLSL
>+
>+} // namespace WebCore
>+
>+#endif
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLTypeNamer.h b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLTypeNamer.h
>new file mode 100644
>index 00000000000..63cd5fe2bc4
>--- /dev/null
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLTypeNamer.h
>@@ -0,0 +1,120 @@
>+/*
>+ * Copyright (C) 2019 Apple Inc. All rights reserved.
>+ *
>+ * Redistribution and use in source and binary forms, with or without
>+ * modification, are permitted provided that the following conditions
>+ * are met:
>+ * 1. Redistributions of source code must retain the above copyright
>+ *    notice, this list of conditions and the following disclaimer.
>+ * 2. Redistributions in binary form must reproduce the above copyright
>+ *    notice, this list of conditions and the following disclaimer in the
>+ *    documentation and/or other materials provided with the distribution.
>+ *
>+ * THIS SOFTWARE IS PROVIDED BY APPLE INC. AND ITS CONTRIBUTORS ``AS IS''
>+ * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO,
>+ * THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
>+ * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL APPLE INC. OR ITS CONTRIBUTORS
>+ * BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
>+ * CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
>+ * SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
>+ * INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
>+ * CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
>+ * ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
>+ * THE POSSIBILITY OF SUCH DAMAGE.
>+ */
>+
>+#pragma once
>+
>+#if ENABLE(WEBGPU)
>+
>+#include "WHLSLVisitor.h"
>+#include <wtf/HashMap.h>
>+#include <wtf/text/StringConcatenate.h>
>+#include <wtf/text/StringConcatenateNumbers.h>
>+#include <wtf/text/WTFString.h>
>+
>+namespace WebCore {
>+
>+namespace WHLSL {
>+
>+namespace AST {
>+
>+class EnumerationDefinition;
>+class EnumerationMember;
>+class NamedType;
>+class NativeTypeDeclaration;
>+class StructureDefinition;
>+class TypeDefinition;
>+class UnnamedType;
>+
>+}
>+
>+class Program;
>+
>+namespace Metal {
>+
>+class BaseTypeNameNode;
>+
>+class TypeNamer : private Visitor {
>+public:
>+    TypeNamer(Program&);
>+    virtual ~TypeNamer();
>+
>+    String metalTypes();
>+
>+    // Must be called after calling metalTypes().
>+    String mangledNameForType(AST::NativeTypeDeclaration&);
>+    String mangledNameForType(AST::UnnamedType&);
>+    String mangledNameForType(AST::NamedType&);
>+    String mangledNameForEnumerationMember(AST::EnumerationMember&);
>+    String mangledNameForStructureElement(AST::StructureElement&);
>+
>+    String generateNextTypeName()
>+    {
>+        return makeString("type", m_typeCount++);
>+    }
>+
>+    String generateNextStructureElementName()
>+    {
>+        return makeString("structureElement", m_structureElementCount++);
>+    }
>+
>+private:
>+    void visit(AST::UnnamedType&) override;
>+    void visit(AST::EnumerationDefinition&) override;
>+    void visit(AST::NativeTypeDeclaration&) override;
>+    void visit(AST::StructureDefinition&) override;
>+    void visit(AST::TypeDefinition&) override;
>+
>+    String generateNextEnumerationMemberName()
>+    {
>+        return makeString("enumerationMember", m_enumerationMemberCount++);
>+    }
>+
>+    void emitNamedTypeDefinition(AST::NamedType&, HashSet<AST::NamedType*>& emittedNamedTypes, HashSet<BaseTypeNameNode*>& emittedUnnamedTypes, StringBuilder&);
>+    void emitUnnamedTypeDefinition(BaseTypeNameNode&, HashSet<AST::NamedType*>& emittedNamedTypes, HashSet<BaseTypeNameNode*>& emittedUnnamedTypes, StringBuilder&);
>+    String metalTypeDeclarations();
>+    String metalTypeDefinitions();
>+
>+    UniqueRef<BaseTypeNameNode> createNameNode(AST::UnnamedType&, BaseTypeNameNode* parent);
>+    size_t insert(AST::UnnamedType&, Vector<UniqueRef<BaseTypeNameNode>>&);
>+
>+    Program& m_program;
>+    Vector<UniqueRef<BaseTypeNameNode>> m_trie;
>+    HashMap<AST::UnnamedType*, BaseTypeNameNode*> m_unnamedTypeMapping;
>+    HashMap<AST::NamedType*, String> m_namedTypeMapping;
>+    HashMap<AST::NamedType*, Vector<std::reference_wrapper<BaseTypeNameNode>>> m_dependencyGraph;
>+    HashMap<AST::EnumerationMember*, String> m_enumerationMemberMapping;
>+    HashMap<AST::StructureElement*, String> m_structureElementMapping;
>+    unsigned m_typeCount { 0 };
>+    unsigned m_enumerationMemberCount { 0 };
>+    unsigned m_structureElementCount { 0 };
>+};
>+
>+} // namespace Metal
>+
>+} // namespace WHLSL
>+
>+} // namespace WebCore
>+
>+#endif
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLChecker.cpp b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLChecker.cpp
>index 60b0efeecac..294dc98be30 100644
>--- a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLChecker.cpp
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLChecker.cpp
>@@ -221,8 +221,8 @@ static bool checkSemantics(Vector<EntryPointItem>& inputItems, Vector<EntryPoint
>         auto checkSemanticTypes = [&](const Vector<EntryPointItem>& items) -> bool {
>             for (auto& item : items) {
>                 auto acceptable = WTF::visit(WTF::makeVisitor([&](const AST::BaseSemantic& semantic) -> bool {
>-                    return semantic.isAcceptableType(item.unnamedType, intrinsics);
>-                }), item.semantic);
>+                    return semantic.isAcceptableType(*item.unnamedType, intrinsics);
>+                }), *item.semantic);
>                 if (!acceptable)
>                     return false;
>             }
>@@ -239,7 +239,7 @@ static bool checkSemantics(Vector<EntryPointItem>& inputItems, Vector<EntryPoint
>             for (auto& item : items) {
>                 auto acceptable = WTF::visit(WTF::makeVisitor([&](const AST::BaseSemantic& semantic) -> bool {
>                     return semantic.isAcceptableForShaderItemDirection(direction, entryPointType);
>-                }), item.semantic);
>+                }), *item.semantic);
>                 if (!acceptable)
>                     return false;
>             }
>@@ -256,11 +256,11 @@ static bool checkSemantics(Vector<EntryPointItem>& inputItems, Vector<EntryPoint
>             for (auto& item : items) {
>                 PODChecker podChecker;
>                 if (is<AST::PointerType>(item.unnamedType))
>-                    podChecker.checkErrorAndVisit(downcast<AST::PointerType>(item.unnamedType).elementType());
>+                    podChecker.checkErrorAndVisit(downcast<AST::PointerType>(*item.unnamedType).elementType());
>                 else if (is<AST::ArrayReferenceType>(item.unnamedType))
>-                    podChecker.checkErrorAndVisit(downcast<AST::ArrayReferenceType>(item.unnamedType).elementType());
>+                    podChecker.checkErrorAndVisit(downcast<AST::ArrayReferenceType>(*item.unnamedType).elementType());
>                 else if (is<AST::ArrayType>(item.unnamedType))
>-                    podChecker.checkErrorAndVisit(downcast<AST::ArrayType>(item.unnamedType).type());
>+                    podChecker.checkErrorAndVisit(downcast<AST::ArrayType>(*item.unnamedType).type());
>                 else
>                     continue;
>                 if (podChecker.error())
>@@ -612,16 +612,12 @@ void Checker::visit(AST::EnumerationDefinition& enumerationDefinition)
>     auto* baseType = ([&]() -> AST::NativeTypeDeclaration* {
>         checkErrorAndVisit(enumerationDefinition.type());
>         auto& baseType = enumerationDefinition.type().unifyNode();
>-        if (!is<AST::UnnamedType>(baseType))
>+        if (!is<AST::NamedType>(baseType))
>             return nullptr;
>-        auto& unnamedBase = downcast<AST::UnnamedType>(baseType);
>-        if (!is<AST::TypeReference>(unnamedBase))
>+        auto& namedType = downcast<AST::NamedType>(baseType);
>+        if (!is<AST::NativeTypeDeclaration>(namedType))
>             return nullptr;
>-        auto& typeReferenceBase = downcast<AST::TypeReference>(unnamedBase);
>-        ASSERT(typeReferenceBase.resolvedType());
>-        if (!is<AST::NativeTypeDeclaration>(*typeReferenceBase.resolvedType()))
>-            return nullptr;
>-        auto& nativeTypeDeclaration = downcast<AST::NativeTypeDeclaration>(*typeReferenceBase.resolvedType());
>+        auto& nativeTypeDeclaration = downcast<AST::NativeTypeDeclaration>(namedType);
>         if (!nativeTypeDeclaration.isInt())
>             return nullptr;
>         return &nativeTypeDeclaration;
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLGatherEntryPointItems.cpp b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLGatherEntryPointItems.cpp
>index a18eac38ea3..45a62a6d288 100644
>--- a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLGatherEntryPointItems.cpp
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLGatherEntryPointItems.cpp
>@@ -71,7 +71,7 @@ public:
>             setError();
>             return;
>         }
>-        m_entryPointItems.append(EntryPointItem(m_typeReferences.last().get(), *m_currentSemantic));
>+        m_entryPointItems.append(EntryPointItem(m_typeReferences.last().get(), *m_currentSemantic, m_path));
>     }
> 
>     void visit(AST::NativeTypeDeclaration& nativeTypeDeclaration)
>@@ -85,7 +85,7 @@ public:
>             return;
>         }
> 
>-        m_entryPointItems.append(EntryPointItem(m_typeReferences.last().get(), *m_currentSemantic));
>+        m_entryPointItems.append(EntryPointItem(m_typeReferences.last().get(), *m_currentSemantic, m_path));
>     }
> 
>     void visit(AST::StructureDefinition& structureDefinition)
>@@ -98,7 +98,9 @@ public:
>         for (auto& structureElement : structureDefinition.structureElements()) {
>             if (structureElement.semantic())
>                 m_currentSemantic = &*structureElement.semantic();
>+            m_path.append(structureElement.name());
>             checkErrorAndVisit(structureElement);
>+            m_path.takeLast();
>         }
>     }
> 
>@@ -122,7 +124,7 @@ public:
>             setError();
>             return;
>         }
>-        m_entryPointItems.append(EntryPointItem(pointerType, *m_currentSemantic));
>+        m_entryPointItems.append(EntryPointItem(pointerType, *m_currentSemantic, m_path));
>     }
> 
>     void visit(AST::ArrayReferenceType& arrayReferenceType)
>@@ -131,7 +133,7 @@ public:
>             setError();
>             return;
>         }
>-        m_entryPointItems.append(EntryPointItem(arrayReferenceType, *m_currentSemantic));
>+        m_entryPointItems.append(EntryPointItem(arrayReferenceType, *m_currentSemantic, m_path));
>     }
> 
>     void visit(AST::ArrayType& arrayType)
>@@ -140,7 +142,7 @@ public:
>             setError();
>             return;
>         }
>-        m_entryPointItems.append(EntryPointItem(arrayType, *m_currentSemantic));
>+        m_entryPointItems.append(EntryPointItem(arrayType, *m_currentSemantic, m_path));
>     }
> 
>     void visit(AST::VariableDeclaration& variableDeclaration)
>@@ -148,11 +150,14 @@ public:
>         ASSERT(!m_currentSemantic);
>         if (variableDeclaration.semantic())
>             m_currentSemantic = &*variableDeclaration.semantic();
>-        if (variableDeclaration.type())
>-            checkErrorAndVisit(*variableDeclaration.type());
>+        ASSERT(variableDeclaration.type());
>+        m_path.append(variableDeclaration.name());
>+        checkErrorAndVisit(*variableDeclaration.type());
>+        m_path.takeLast();
>     }
> 
> private:
>+    Vector<String> m_path;
>     const Intrinsics& m_intrinsics;
>     AST::Semantic* m_currentSemantic { nullptr };
>     Vector<std::reference_wrapper<AST::TypeReference>> m_typeReferences;
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLGatherEntryPointItems.h b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLGatherEntryPointItems.h
>index 20367be3e29..6664e2e342e 100644
>--- a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLGatherEntryPointItems.h
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLGatherEntryPointItems.h
>@@ -30,6 +30,7 @@
> #include "WHLSLSemantic.h"
> #include <wtf/Optional.h>
> #include <wtf/Vector.h>
>+#include <wtf/text/WTFString.h>
> 
> namespace WebCore {
> 
>@@ -45,14 +46,16 @@ class Intrinsics;
> class UnnamedType;
> 
> struct EntryPointItem {
>-    EntryPointItem(AST::UnnamedType& unnamedType, AST::Semantic& semantic)
>-        : unnamedType(unnamedType)
>-        , semantic(semantic)
>+    EntryPointItem(AST::UnnamedType& unnamedType, AST::Semantic& semantic, Vector<String>& path)
>+        : unnamedType(&unnamedType)
>+        , semantic(&semantic)
>+        , path(path)
>     {
>     }
> 
>-    AST::UnnamedType& unnamedType;
>-    AST::Semantic& semantic;
>+    AST::UnnamedType* unnamedType;
>+    AST::Semantic* semantic;
>+    Vector<String> path;
> };
> 
> struct EntryPointItems {
>diff --git a/Source/WebCore/Sources.txt b/Source/WebCore/Sources.txt
>index dc2e4bf0dd7..649dd7e0d61 100644
>--- a/Source/WebCore/Sources.txt
>+++ b/Source/WebCore/Sources.txt
>@@ -326,6 +326,11 @@ Modules/webgpu/WHLSL/WHLSLRecursionChecker.cpp
> Modules/webgpu/WHLSL/WHLSLVisitor.cpp
> Modules/webgpu/WHLSL/WHLSLLiteralTypeChecker.cpp
> Modules/webgpu/WHLSL/WHLSLHighZombieFinder.cpp
>+Modules/webgpu/WHLSL/Metal/WHLSLTypeNamer.cpp
>+Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.cpp
>+Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.cpp
>+Modules/webgpu/WHLSL/Metal/WHLSLNativeFunctionWriter.cpp
>+Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp
> Modules/webgpu/WHLSL/WHLSLFunctionStageChecker.cpp
> Modules/webgpu/WHLSL/AST/WHLSLTypeArgument.cpp
> Modules/webgpu/WHLSL/AST/WHLSLBuiltInSemantic.cpp
>diff --git a/Source/WebCore/WebCore.xcodeproj/project.pbxproj b/Source/WebCore/WebCore.xcodeproj/project.pbxproj
>index 5c9ab0560bf..0ec5af5dfb5 100644
>--- a/Source/WebCore/WebCore.xcodeproj/project.pbxproj
>+++ b/Source/WebCore/WebCore.xcodeproj/project.pbxproj
>@@ -6460,6 +6460,16 @@
> 		1CDD45E60BA9C84600F90147 /* Base.xcconfig */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = text.xcconfig; path = Base.xcconfig; sourceTree = "<group>"; };
> 		1CECB3A821F2B67300F44542 /* WHLSLStatementBehaviorChecker.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = WHLSLStatementBehaviorChecker.h; sourceTree = "<group>"; };
> 		1CECB3A921F2B67300F44542 /* WHLSLStatementBehaviorChecker.cpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.cpp; path = WHLSLStatementBehaviorChecker.cpp; sourceTree = "<group>"; };
>+		1CECB3AF21F2B98400F44542 /* WHLSLFunctionWriter.cpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.cpp; path = WHLSLFunctionWriter.cpp; sourceTree = "<group>"; };
>+		1CECB3B021F2B98500F44542 /* WHLSLTypeNamer.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = WHLSLTypeNamer.h; sourceTree = "<group>"; };
>+		1CECB3B121F2B98500F44542 /* WHLSLTypeNamer.cpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.cpp; path = WHLSLTypeNamer.cpp; sourceTree = "<group>"; };
>+		1CECB3B221F2B98600F44542 /* WHLSLFunctionWriter.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = WHLSLFunctionWriter.h; sourceTree = "<group>"; };
>+		1CECB3B521F50AC700F44542 /* WHLSLMetalCodeGenerator.cpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.cpp; path = WHLSLMetalCodeGenerator.cpp; sourceTree = "<group>"; };
>+		1CECB3B621F50AC700F44542 /* WHLSLMetalCodeGenerator.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = WHLSLMetalCodeGenerator.h; sourceTree = "<group>"; };
>+		1CECB3B821F50D1000F44542 /* WHLSLNativeFunctionWriter.cpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.cpp; path = WHLSLNativeFunctionWriter.cpp; sourceTree = "<group>"; };
>+		1CECB3B921F50D1000F44542 /* WHLSLNativeFunctionWriter.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = WHLSLNativeFunctionWriter.h; sourceTree = "<group>"; };
>+		1CECB3BA21F511AA00F44542 /* WHLSLEntryPointScaffolding.cpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.cpp; path = WHLSLEntryPointScaffolding.cpp; sourceTree = "<group>"; };
>+		1CECB3BB21F511AA00F44542 /* WHLSLEntryPointScaffolding.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = WHLSLEntryPointScaffolding.h; sourceTree = "<group>"; };
> 		1CFAE3220A6D6A3F0032593D /* libobjc.dylib */ = {isa = PBXFileReference; lastKnownFileType = "compiled.mach-o.dylib"; name = libobjc.dylib; path = /usr/lib/libobjc.dylib; sourceTree = "<absolute>"; };
> 		1DC553FD211BA12A004B780E /* NavigatorShare.idl */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = text; path = NavigatorShare.idl; sourceTree = "<group>"; };
> 		1DC553FF211BA841004B780E /* ShareData.idl */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = text; path = ShareData.idl; sourceTree = "<group>"; };
>@@ -17150,6 +17160,23 @@
> 			tabWidth = 4;
> 			usesTabs = 0;
> 		};
>+		1CECB3AD21F2B96400F44542 /* Metal */ = {
>+			isa = PBXGroup;
>+			children = (
>+				1CECB3BA21F511AA00F44542 /* WHLSLEntryPointScaffolding.cpp */,
>+				1CECB3BB21F511AA00F44542 /* WHLSLEntryPointScaffolding.h */,
>+				1CECB3AF21F2B98400F44542 /* WHLSLFunctionWriter.cpp */,
>+				1CECB3B221F2B98600F44542 /* WHLSLFunctionWriter.h */,
>+				1CECB3B521F50AC700F44542 /* WHLSLMetalCodeGenerator.cpp */,
>+				1CECB3B621F50AC700F44542 /* WHLSLMetalCodeGenerator.h */,
>+				1CECB3B821F50D1000F44542 /* WHLSLNativeFunctionWriter.cpp */,
>+				1CECB3B921F50D1000F44542 /* WHLSLNativeFunctionWriter.h */,
>+				1CECB3B121F2B98500F44542 /* WHLSLTypeNamer.cpp */,
>+				1CECB3B021F2B98500F44542 /* WHLSLTypeNamer.h */,
>+			);
>+			path = Metal;
>+			sourceTree = "<group>";
>+		};
> 		26B9998D1803ADFA00D01121 /* cssjit */ = {
> 			isa = PBXGroup;
> 			children = (
>@@ -25562,6 +25589,7 @@
> 			isa = PBXGroup;
> 			children = (
> 				1CA0C2F621EEDAD200A11860 /* AST */,
>+				1CECB3AD21F2B96400F44542 /* Metal */,
> 				C234A9B221E92C1F003C984D /* WHLSLCheckDuplicateFunctions.cpp */,
> 				C234A9AE21E92C1A003C984D /* WHLSLCheckDuplicateFunctions.h */,
> 				1C840B9B21EC400900D0500D /* WHLSLChecker.cpp */,

Actions: View | Formatted Diff | Diff

Attachments on bug 193531: 359362 | 359401 | 359459 | 359463 | 359473 | 359643 | 359644 | 359646 | 359647