[WHLSL] Implement Metal code generation
authormmaxfield@apple.com <mmaxfield@apple.com@268f45cc-cd09-0410-ab3c-d52691b4dbfc>
Mon, 21 Jan 2019 07:31:29 +0000 (07:31 +0000)
committermmaxfield@apple.com <mmaxfield@apple.com@268f45cc-cd09-0410-ab3c-d52691b4dbfc>
Mon, 21 Jan 2019 07:31:29 +0000 (07:31 +0000)
https://bugs.webkit.org/show_bug.cgi?id=193531

Reviewed by Dean Jackson.

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 its 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: Renamed from Source/WebCore/Modules/webgpu/WHLSL/WHLSLLoopChecker.h.
* Sources.txt:
* WebCore.xcodeproj/project.pbxproj:

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

26 files changed:
Source/WebCore/ChangeLog
Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLBuiltInSemantic.h
Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLExpression.h
Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLFunctionDeclaration.h
Source/WebCore/Modules/webgpu/WHLSL/AST/WHLSLStructureDefinition.h
Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp [new file with mode: 0644]
Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.h [new file with mode: 0644]
Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.cpp [new file with mode: 0644]
Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLFunctionWriter.h [new file with mode: 0644]
Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.cpp [new file with mode: 0644]
Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLMetalCodeGenerator.h [new file with mode: 0644]
Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLNativeFunctionWriter.cpp [new file with mode: 0644]
Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLNativeFunctionWriter.h [new file with mode: 0644]
Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLNativeTypeWriter.cpp [new file with mode: 0644]
Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLNativeTypeWriter.h [new file with mode: 0644]
Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLTypeNamer.cpp [new file with mode: 0644]
Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLTypeNamer.h [new file with mode: 0644]
Source/WebCore/Modules/webgpu/WHLSL/WHLSLChecker.cpp
Source/WebCore/Modules/webgpu/WHLSL/WHLSLGatherEntryPointItems.cpp
Source/WebCore/Modules/webgpu/WHLSL/WHLSLGatherEntryPointItems.h
Source/WebCore/Modules/webgpu/WHLSL/WHLSLParser.cpp
Source/WebCore/Modules/webgpu/WHLSL/WHLSLStatementBehaviorChecker.cpp
Source/WebCore/Modules/webgpu/WHLSL/WHLSLSynthesizeConstructors.cpp
Source/WebCore/Modules/websockets/WebSocketChannel.cpp
Source/WebCore/Sources.txt
Source/WebCore/WebCore.xcodeproj/project.pbxproj

index a001fbb..9cf7f91 100644 (file)
@@ -1,5 +1,142 @@
 2019-01-20  Myles C. Maxfield  <mmaxfield@apple.com>
 
+        [WHLSL] Implement Metal code generation
+        https://bugs.webkit.org/show_bug.cgi?id=193531
+
+        Reviewed by Dean Jackson.
+
+        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 its 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: Renamed from Source/WebCore/Modules/webgpu/WHLSL/WHLSLLoopChecker.h.
+        * Sources.txt:
+        * WebCore.xcodeproj/project.pbxproj:
+
+2019-01-20  Myles C. Maxfield  <mmaxfield@apple.com>
+
         [WHLSL] Add the statement behavior checker
         https://bugs.webkit.org/show_bug.cgi?id=193487
 
index 43b6986..7b36caa 100644 (file)
@@ -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
     {
index 8fef629..01f2ffb 100644 (file)
@@ -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)
     {
index 5082858..c15466a 100644 (file)
@@ -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; }
index 456c6e1..0260d45 100644 (file)
@@ -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 (file)
index 0000000..48fcf23
--- /dev/null
@@ -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&, 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 (file)
index 0000000..8d65fad
--- /dev/null
@@ -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& existingVariableName, 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 (file)
index 0000000..019a267
--- /dev/null
@@ -0,0 +1,632 @@
+/*
+ * 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());
+            ASSERT(m_stack.isEmpty());
+            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());
+            ASSERT(m_stack.isEmpty());
+            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 nothing 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_stack.takeLast();
+        }
+        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(m_stack.takeLast(), 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");
+        checkErrorAndVisit(switchCase.block());
+        // FIXME: Figure out whether we need to break or fallthrough.
+        CRASH();
+    }
+
+    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 (file)
index 0000000..9960e5e
--- /dev/null
@@ -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 (file)
index 0000000..10c0530
--- /dev/null
@@ -0,0 +1,66 @@
+/*
+ * 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("#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 (file)
index 0000000..e088e01
--- /dev/null
@@ -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 (file)
index 0000000..b4d3706
--- /dev/null
@@ -0,0 +1,399 @@
+/*
+ * 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 "WHLSLPointerType.h"
+#include "WHLSLTypeNamer.h"
+#include "WHLSLUnnamedType.h"
+#include "WHLSLVariableDeclaration.h"
+#include <wtf/text/StringBuilder.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;
+    }
+}
+
+static String atomicName(String input)
+{
+    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;
+}
+
+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("=")) {
+            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 = atomicName(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 (file)
index 0000000..2e9bd4e
--- /dev/null
@@ -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/WHLSLNativeTypeWriter.cpp b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLNativeTypeWriter.cpp
new file mode 100644 (file)
index 0000000..a7c9bf5
--- /dev/null
@@ -0,0 +1,256 @@
+/*
+ * 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 "WHLSLNativeTypeWriter.h"
+
+#if ENABLE(WEBGPU)
+
+#include "WHLSLNamedType.h"
+#include "WHLSLNativeTypeDeclaration.h"
+#include "WHLSLTypeReference.h"
+
+namespace WebCore {
+
+namespace WHLSL {
+
+namespace Metal {
+
+String writeNativeType(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, '>');
+}
+
+} // namespace Metal
+
+} // namespace WHLSL
+
+} // namespace WebCore
+
+#endif
diff --git a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLNativeTypeWriter.h b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLNativeTypeWriter.h
new file mode 100644 (file)
index 0000000..9430b0b
--- /dev/null
@@ -0,0 +1,52 @@
+/*
+ * 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 NativeTypeDeclaration;
+
+}
+
+namespace Metal {
+
+String writeNativeType(AST::NativeTypeDeclaration&);
+
+} // namespace Metal
+
+} // namespace WHLSL
+
+} // namespace WebCore
+
+#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 (file)
index 0000000..6ff36f9
--- /dev/null
@@ -0,0 +1,503 @@
+/*
+ * 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 "WHLSLNativeTypeWriter.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)
+{
+    return writeNativeType(nativeTypeDeclaration);
+}
+
+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 (file)
index 0000000..63cd5fe
--- /dev/null
@@ -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
index 60b0efe..294dc98 100644 (file)
@@ -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;
index a18eac3..45a62a6 100644 (file)
@@ -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;
index 20367be..6664e2e 100644 (file)
@@ -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 {
index e58b5a5..baa6caf 100644 (file)
@@ -245,7 +245,7 @@ static Expected<unsigned, Parser::Error> uintLiteralToUint(StringView text)
         text = text.substring(2);
         base = 16;
     }
-    ASSERT(text.endsWith("u"_str));
+    ASSERT(text.endsWith("u"));
     text = text.substring(0, text.length() - 1);
     unsigned result = 0;
     for (auto codePoint : text.codePoints()) {
index eb8f88e..b4bcd61 100644 (file)
 
 #if ENABLE(WEBGPU)
 
+#include "WHLSLBlock.h"
+#include "WHLSLBreak.h"
+#include "WHLSLContinue.h"
 #include "WHLSLDoWhileLoop.h"
+#include "WHLSLEffectfulExpressionStatement.h"
+#include "WHLSLFallthrough.h"
 #include "WHLSLForLoop.h"
 #include "WHLSLIfStatement.h"
 #include "WHLSLInferTypes.h"
 #include "WHLSLProgram.h"
+#include "WHLSLReturn.h"
+#include "WHLSLSwitchCase.h"
 #include "WHLSLSwitchStatement.h"
+#include "WHLSLTrap.h"
+#include "WHLSLVariableDeclarationsStatement.h"
+#include "WHLSLVisitor.h"
 #include "WHLSLWhileLoop.h"
 #include <cstdint>
 #include <wtf/OptionSet.h>
index 86f5cdb..ed32b64 100644 (file)
 
 #if ENABLE(WEBGPU)
 
+#include "WHLSLArrayType.h"
+#include "WHLSLArrayReferenceType.h"
+#include "WHLSLEnumerationDefinition.h"
 #include "WHLSLNativeFunctionDeclaration.h"
+#include "WHLSLNativeTypeDeclaration.h"
+#include "WHLSLPointerType.h"
 #include "WHLSLProgram.h"
+#include "WHLSLStructureDefinition.h"
 #include "WHLSLTypeReference.h"
 #include "WHLSLVariableDeclaration.h"
 #include "WHLSLVisitor.h"
index 4ea3b12..24aea60 100644 (file)
@@ -38,6 +38,7 @@
 #include "FileError.h"
 #include "FileReaderLoader.h"
 #include "Frame.h"
+#include "FrameLoader.h"
 #include "InspectorInstrumentation.h"
 #include "Logging.h"
 #include "NetworkingContext.h"
index dc2e4bf..274b088 100644 (file)
@@ -326,6 +326,12 @@ 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/Metal/WHLSLNativeTypeWriter.cpp
 Modules/webgpu/WHLSL/WHLSLFunctionStageChecker.cpp
 Modules/webgpu/WHLSL/AST/WHLSLTypeArgument.cpp
 Modules/webgpu/WHLSL/AST/WHLSLBuiltInSemantic.cpp
index 5c9ab05..b5dee9c 100644 (file)
                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>"; };
+               1CECB3C621F59C8700F44542 /* WHLSLNativeTypeWriter.cpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.cpp; path = WHLSLNativeTypeWriter.cpp; sourceTree = "<group>"; };
+               1CECB3C721F59C8700F44542 /* WHLSLNativeTypeWriter.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = WHLSLNativeTypeWriter.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>"; };
                        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 */,
+                               1CECB3C621F59C8700F44542 /* WHLSLNativeTypeWriter.cpp */,
+                               1CECB3C721F59C8700F44542 /* WHLSLNativeTypeWriter.h */,
+                       );
+                       path = Metal;
+                       sourceTree = "<group>";
+               };
                26B9998D1803ADFA00D01121 /* cssjit */ = {
                        isa = PBXGroup;
                        children = (
                        isa = PBXGroup;
                        children = (
                                1CA0C2F621EEDAD200A11860 /* AST */,
+                               1CECB3AD21F2B96400F44542 /* Metal */,
                                C234A9B221E92C1F003C984D /* WHLSLCheckDuplicateFunctions.cpp */,
                                C234A9AE21E92C1A003C984D /* WHLSLCheckDuplicateFunctions.h */,
                                1C840B9B21EC400900D0500D /* WHLSLChecker.cpp */,