Unreviewed, rolling out r246396 and r246397.
authorsroberts@apple.com <sroberts@apple.com@268f45cc-cd09-0410-ab3c-d52691b4dbfc>
Thu, 13 Jun 2019 19:06:31 +0000 (19:06 +0000)
committersroberts@apple.com <sroberts@apple.com@268f45cc-cd09-0410-ab3c-d52691b4dbfc>
Thu, 13 Jun 2019 19:06:31 +0000 (19:06 +0000)
https://bugs.webkit.org/show_bug.cgi?id=198837

Breaks internal builds (Requested by ShawnRoberts on #webkit).

Reverted changesets:

"[WHLSL] Hook up compute"
https://bugs.webkit.org/show_bug.cgi?id=198644
https://trac.webkit.org/changeset/246396

"[WHLSL] Hook up compute"
https://bugs.webkit.org/show_bug.cgi?id=198644
https://trac.webkit.org/changeset/246397

Patch by Commit Queue <commit-queue@webkit.org> on 2019-06-13

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

24 files changed:
LayoutTests/ChangeLog
LayoutTests/webgpu/compute-squares-expected.txt [new file with mode: 0644]
LayoutTests/webgpu/compute-squares.html [new file with mode: 0644]
LayoutTests/webgpu/whlsl-compute-expected.txt [deleted file]
LayoutTests/webgpu/whlsl-compute.html [deleted file]
Source/WebCore/ChangeLog
Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp
Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.h
Source/WebCore/Modules/webgpu/WHLSL/WHLSLComputeDimensions.cpp [deleted file]
Source/WebCore/Modules/webgpu/WHLSL/WHLSLComputeDimensions.h [deleted file]
Source/WebCore/Modules/webgpu/WHLSL/WHLSLGatherEntryPointItems.cpp
Source/WebCore/Modules/webgpu/WHLSL/WHLSLPrepare.cpp
Source/WebCore/Modules/webgpu/WHLSL/WHLSLPrepare.h
Source/WebCore/Modules/webgpu/WHLSL/WHLSLPropertyResolver.cpp
Source/WebCore/Modules/webgpu/WHLSL/WHLSLStandardLibrary.txt
Source/WebCore/Sources.txt
Source/WebCore/SourcesCocoa.txt
Source/WebCore/WebCore.xcodeproj/project.pbxproj
Source/WebCore/platform/graphics/gpu/GPUComputePipeline.h
Source/WebCore/platform/graphics/gpu/cocoa/GPUComputePassEncoderMetal.mm
Source/WebCore/platform/graphics/gpu/cocoa/GPUComputePipelineMetal.mm
Source/WebCore/platform/graphics/gpu/cocoa/GPUPipelineMetalConvertLayout.cpp [deleted file]
Source/WebCore/platform/graphics/gpu/cocoa/GPUPipelineMetalConvertLayout.h [deleted file]
Source/WebCore/platform/graphics/gpu/cocoa/GPURenderPipelineMetal.mm

index 5d4d28b..9ba3964 100644 (file)
@@ -1,3 +1,20 @@
+2019-06-13  Commit Queue  <commit-queue@webkit.org>
+
+        Unreviewed, rolling out r246396 and r246397.
+        https://bugs.webkit.org/show_bug.cgi?id=198837
+
+        Breaks internal builds (Requested by ShawnRoberts on #webkit).
+
+        Reverted changesets:
+
+        "[WHLSL] Hook up compute"
+        https://bugs.webkit.org/show_bug.cgi?id=198644
+        https://trac.webkit.org/changeset/246396
+
+        "[WHLSL] Hook up compute"
+        https://bugs.webkit.org/show_bug.cgi?id=198644
+        https://trac.webkit.org/changeset/246397
+
 2019-06-13  Antti Koivisto  <antti@apple.com>
 
         twitch.tv: embedded video hovers down the screen when scrolling on iPad
diff --git a/LayoutTests/webgpu/compute-squares-expected.txt b/LayoutTests/webgpu/compute-squares-expected.txt
new file mode 100644 (file)
index 0000000..8b13789
--- /dev/null
@@ -0,0 +1 @@
+
diff --git a/LayoutTests/webgpu/compute-squares.html b/LayoutTests/webgpu/compute-squares.html
new file mode 100644 (file)
index 0000000..39bffc0
--- /dev/null
@@ -0,0 +1,78 @@
+<!DOCTYPE html><!-- webkit-test-runner [ experimental:WebGPUEnabled=true ] -->
+<meta charset=utf-8>
+<title>Execute a simple compute shader with Web GPU.</title>
+<body>
+<script src="../resources/testharness.js"></script>
+<script src="../resources/testharnessreport.js"></script>
+<script src="js/webgpu-functions.js"></script>
+<script>
+if (window.testRunner)
+    testRunner.waitUntilDone();
+
+const data = new Uint32Array([2, 3, 4, 5, 6, 7, 8, 9, 10]);
+
+const dataBinding = 0;
+const bindGroupIndex = 0;
+
+const shaderCode = `
+#include <metal_stdlib>
+
+struct Data {
+    device unsigned* numbers [[id(${dataBinding})]];
+};
+
+kernel void compute(device Data& data [[buffer(${bindGroupIndex})]], unsigned gid [[thread_position_in_grid]])
+{
+    if (gid >= ${data.length})
+        return;
+
+    unsigned original = data.numbers[gid];
+    data.numbers[gid] = original * original;
+}
+`
+
+promise_test(async () => {
+    
+    const device = await getBasicDevice();
+    
+    const shaderModule = device.createShaderModule({ code: shaderCode, isWHLSL: false });
+    const computeStageDescriptor = { module: shaderModule, entryPoint: "compute" };
+    const pipeline = device.createComputePipeline({ computeStage: computeStageDescriptor });
+    
+    const dataBuffer = createBufferWithData(device, { size: data.byteLength, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.MAP_READ }, data.buffer);
+    
+    const bgLayoutBinding = { binding: dataBinding, visibility: GPUShaderStageBit.COMPUTE, type: "storage-buffer" };
+    const bgLayout = device.createBindGroupLayout({ bindings: [bgLayoutBinding] });
+    
+    const bufferBinding = { buffer: dataBuffer, size: data.byteLength };
+    const bgBinding = { binding: dataBinding, resource: bufferBinding };
+    
+    const bindGroupDescriptor = { layout: bgLayout, bindings: [bgBinding] };
+    const bindGroup = device.createBindGroup(bindGroupDescriptor);
+    
+    const commandEncoder = device.createCommandEncoder();
+    const passEncoder = commandEncoder.beginComputePass();
+    
+    passEncoder.setBindGroup(bindGroupIndex, bindGroup);
+    
+    passEncoder.setPipeline(pipeline);
+    
+    // One thread group.
+    passEncoder.dispatch(1, 1, 1);
+    passEncoder.endPass();
+    
+    device.getQueue().submit([commandEncoder.finish()]);
+    
+    const readDataArrayBuffer = await dataBuffer.mapReadAsync();
+    assert_not_equals(readDataArrayBuffer, null, "Async read promise resolved successfully");
+    
+    const readData = new Uint32Array(readDataArrayBuffer);
+
+    for (var i = 0; i < readData.length; ++i)
+        assert_equals(readData[i], data[i] * data[i], "Data was succesfully squared");
+
+    if (window.testRunner)
+        testRunner.notifyDone();
+}, "Successfully executed a basic compute pass");
+</script>
+</body>
\ No newline at end of file
diff --git a/LayoutTests/webgpu/whlsl-compute-expected.txt b/LayoutTests/webgpu/whlsl-compute-expected.txt
deleted file mode 100644 (file)
index 6e1668a..0000000
+++ /dev/null
@@ -1,12 +0,0 @@
-PASS successfullyParsed is true
-
-TEST COMPLETE
-PASS resultsFloat32Array[0] is 2
-PASS resultsFloat32Array[1] is 4
-PASS resultsFloat32Array[2] is 6
-PASS resultsFloat32Array[3] is 8
-PASS resultsFloat32Array[4] is 5
-PASS resultsFloat32Array[5] is 6
-PASS resultsFloat32Array[6] is 7
-PASS resultsFloat32Array[7] is 8
-
diff --git a/LayoutTests/webgpu/whlsl-compute.html b/LayoutTests/webgpu/whlsl-compute.html
deleted file mode 100644 (file)
index 9058407..0000000
+++ /dev/null
@@ -1,88 +0,0 @@
-<!DOCTYPE html>
-<html>
-<head>
-<script src="../resources/js-test-pre.js"></script>
-</head>
-<body>
-<script>
-const shaderSource = `
-[numthreads(2, 1, 1)]
-compute void computeShader(device float[] buffer : register(u0), float3 threadID : SV_DispatchThreadID) {
-    buffer[uint(threadID.x)] = buffer[uint(threadID.x)] * 2.0;
-}
-`;
-let resultsFloat32Array;
-async function start() {
-    const adapter = await navigator.gpu.requestAdapter();
-    const device = await adapter.requestDevice();
-
-    const shaderModule = device.createShaderModule({code: shaderSource, isWHLSL: true});
-    const computeStage = {module: shaderModule, entryPoint: "computeShader"};
-
-    const bindGroupLayoutDescriptor = {bindings: [{binding: 0, visibility: 7, type: "storage-buffer"}]};
-    const bindGroupLayout = device.createBindGroupLayout(bindGroupLayoutDescriptor);
-    const pipelineLayoutDescriptor = {bindGroupLayouts: [bindGroupLayout]};
-    const pipelineLayout = device.createPipelineLayout(pipelineLayoutDescriptor);
-
-    const computePipelineDescriptor = {computeStage, layout: pipelineLayout};
-    const computePipeline = device.createComputePipeline(computePipelineDescriptor);
-
-    const bufferDescriptor = {size: Float32Array.BYTES_PER_ELEMENT * 8, usage: GPUBufferUsage.MAP_WRITE | GPUBufferUsage.TRANSFER_SRC};
-    const buffer = device.createBuffer(bufferDescriptor);
-    const bufferArrayBuffer = await buffer.mapWriteAsync();
-    const bufferFloat32Array = new Float32Array(bufferArrayBuffer);
-    bufferFloat32Array[0] = 1;
-    bufferFloat32Array[1] = 2;
-    bufferFloat32Array[2] = 3;
-    bufferFloat32Array[3] = 4;
-    bufferFloat32Array[4] = 5;
-    bufferFloat32Array[5] = 6;
-    bufferFloat32Array[6] = 7;
-    bufferFloat32Array[7] = 8;
-    buffer.unmap();
-
-    const resultsBufferDescriptor = {size: Float32Array.BYTES_PER_ELEMENT * 8, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.TRANSFER_DST | GPUBufferUsage.MAP_READ};
-    const resultsBuffer = device.createBuffer(resultsBufferDescriptor);
-
-    const bufferBinding = {buffer: resultsBuffer, size: 4};
-    const bindGroupBinding = {binding: 0, resource: bufferBinding};
-    const bindGroupDescriptor = {layout: bindGroupLayout, bindings: [bindGroupBinding]};
-    const bindGroup = device.createBindGroup(bindGroupDescriptor);
-
-    const commandEncoder = device.createCommandEncoder(); // {}
-    commandEncoder.copyBufferToBuffer(buffer, 0, resultsBuffer, 0, Float32Array.BYTES_PER_ELEMENT * 8);
-    const computePassEncoder = commandEncoder.beginComputePass();
-    computePassEncoder.setPipeline(computePipeline);
-    computePassEncoder.setBindGroup(0, bindGroup);
-    computePassEncoder.dispatch(2, 1, 1);
-    computePassEncoder.endPass();
-    const commandBuffer = commandEncoder.finish();
-    device.getQueue().submit([commandBuffer]);
-
-    const resultsArrayBuffer = await resultsBuffer.mapReadAsync();
-    resultsFloat32Array = new Float32Array(resultsArrayBuffer);
-    shouldBe("resultsFloat32Array[0]", "2");
-    shouldBe("resultsFloat32Array[1]", "4");
-    shouldBe("resultsFloat32Array[2]", "6");
-    shouldBe("resultsFloat32Array[3]", "8");
-    shouldBe("resultsFloat32Array[4]", "5");
-    shouldBe("resultsFloat32Array[5]", "6");
-    shouldBe("resultsFloat32Array[6]", "7");
-    shouldBe("resultsFloat32Array[7]", "8");
-    resultsBuffer.unmap();
-}
-if (window.testRunner)
-    testRunner.waitUntilDone();
-window.addEventListener("load", function() {
-    start().then(function() {
-        if (window.testRunner)
-            testRunner.notifyDone();
-    }, function() {
-        if (window.testRunner)
-            testRunner.notifyDone();
-    });
-});
-</script>
-<script src="../resources/js-test-post.js"></script>
-</body>
-</html>
index e7dc3d9..6ef0baa 100644 (file)
@@ -1,3 +1,20 @@
+2019-06-13  Commit Queue  <commit-queue@webkit.org>
+
+        Unreviewed, rolling out r246396 and r246397.
+        https://bugs.webkit.org/show_bug.cgi?id=198837
+
+        Breaks internal builds (Requested by ShawnRoberts on #webkit).
+
+        Reverted changesets:
+
+        "[WHLSL] Hook up compute"
+        https://bugs.webkit.org/show_bug.cgi?id=198644
+        https://trac.webkit.org/changeset/246396
+
+        "[WHLSL] Hook up compute"
+        https://bugs.webkit.org/show_bug.cgi?id=198644
+        https://trac.webkit.org/changeset/246397
+
 2019-06-13  Antti Koivisto  <antti@apple.com>
 
         twitch.tv: embedded video hovers down the screen when scrolling on iPad
index 3baa37f..da8c057 100644 (file)
@@ -178,41 +178,6 @@ Optional<String> EntryPointScaffolding::resourceSignature()
     return stringBuilder.toString();
 }
 
-static String internalTypeForSemantic(const AST::BuiltInSemantic& builtInSemantic)
-{
-    switch (builtInSemantic.variable()) {
-    case AST::BuiltInSemantic::Variable::SVInstanceID:
-        return "uint"_str;
-    case AST::BuiltInSemantic::Variable::SVVertexID:
-        return "uint"_str;
-    case AST::BuiltInSemantic::Variable::PSize:
-        return "float"_str;
-    case AST::BuiltInSemantic::Variable::SVPosition:
-        return "float4"_str;
-    case AST::BuiltInSemantic::Variable::SVIsFrontFace:
-        return "bool"_str;
-    case AST::BuiltInSemantic::Variable::SVSampleIndex:
-        return "uint"_str;
-    case AST::BuiltInSemantic::Variable::SVInnerCoverage:
-        return "uint"_str;
-    case AST::BuiltInSemantic::Variable::SVTarget:
-        return String();
-    case AST::BuiltInSemantic::Variable::SVDepth:
-        return "float"_str;
-    case AST::BuiltInSemantic::Variable::SVCoverage:
-        return "uint"_str;
-    case AST::BuiltInSemantic::Variable::SVDispatchThreadID:
-        return "uint3"_str;
-    case AST::BuiltInSemantic::Variable::SVGroupID:
-        return "uint3"_str;
-    case AST::BuiltInSemantic::Variable::SVGroupIndex:
-        return "uint"_str;
-    default:
-        ASSERT(builtInSemantic.variable() == AST::BuiltInSemantic::Variable::SVGroupThreadID);
-        return "uint3"_str;
-    }
-}
-
 Optional<String> EntryPointScaffolding::builtInsSignature()
 {
     if (!m_namedBuiltIns.size())
@@ -225,11 +190,9 @@ Optional<String> EntryPointScaffolding::builtInsSignature()
         auto& namedBuiltIn = m_namedBuiltIns[i];
         auto& item = m_entryPointItems.inputs[namedBuiltIn.indexInEntryPointItems];
         auto& builtInSemantic = WTF::get<AST::BuiltInSemantic>(*item.semantic);
-        auto internalType = internalTypeForSemantic(builtInSemantic);
-        if (internalType.isNull())
-            internalType = m_typeNamer.mangledNameForType(*item.unnamedType);
+        auto mangledTypeName = m_typeNamer.mangledNameForType(*item.unnamedType);
         auto variableName = namedBuiltIn.variableName;
-        stringBuilder.append(makeString(internalType, ' ', variableName, ' ', attributeForSemantic(builtInSemantic)));
+        stringBuilder.append(makeString(mangledTypeName, ' ', variableName, ' ', attributeForSemantic(builtInSemantic)));
     }
     return stringBuilder.toString();
 }
@@ -336,11 +299,9 @@ String EntryPointScaffolding::unpackResourcesAndNamedBuiltIns()
     }
 
     for (auto& namedBuiltIn : m_namedBuiltIns) {
-        auto& item = m_entryPointItems.inputs[namedBuiltIn.indexInEntryPointItems];
-        auto& path = item.path;
+        auto& path = m_entryPointItems.inputs[namedBuiltIn.indexInEntryPointItems].path;
         auto& variableName = namedBuiltIn.variableName;
-        auto mangledTypeName = m_typeNamer.mangledNameForType(*item.unnamedType);
-        stringBuilder.append(makeString(mangledInputPath(path), " = ", mangledTypeName, '(', variableName, ");\n"));
+        stringBuilder.append(makeString(mangledInputPath(path), " = ", variableName, ";\n"));
     }
     return stringBuilder.toString();
 }
@@ -363,13 +324,8 @@ VertexEntryPointScaffolding::VertexEntryPointScaffolding(AST::FunctionDefinition
 
     m_namedOutputs.reserveInitialCapacity(m_entryPointItems.outputs.size());
     for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
-        auto& outputItem = m_entryPointItems.outputs[i];
         NamedOutput namedOutput;
         namedOutput.elementName = m_typeNamer.generateNextStructureElementName();
-        if (WTF::holds_alternative<AST::BuiltInSemantic>(*outputItem.semantic))
-            namedOutput.internalTypeName = internalTypeForSemantic(WTF::get<AST::BuiltInSemantic>(*outputItem.semantic));
-        if (namedOutput.internalTypeName.isNull())
-            namedOutput.internalTypeName = m_typeNamer.mangledNameForType(*outputItem.unnamedType);
         m_namedOutputs.uncheckedAppend(WTFMove(namedOutput));
     }
 }
@@ -390,10 +346,10 @@ String VertexEntryPointScaffolding::helperTypes()
     stringBuilder.append(makeString("struct ", m_returnStructName, " {\n"));
     for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
         auto& outputItem = m_entryPointItems.outputs[i];
-        auto& internalTypeName = m_namedOutputs[i].internalTypeName;
+        auto mangledTypeName = m_typeNamer.mangledNameForType(*outputItem.unnamedType);
         auto elementName = m_namedOutputs[i].elementName;
         auto attribute = attributeForSemantic(*outputItem.semantic);
-        stringBuilder.append(makeString("    ", internalTypeName, ' ', elementName, ' ', attribute, ";\n"));
+        stringBuilder.append(makeString("    ", mangledTypeName, ' ', elementName, ' ', attribute, ";\n"));
     }
     stringBuilder.append("};\n\n");
 
@@ -442,9 +398,8 @@ String VertexEntryPointScaffolding::pack(const String& inputVariableName, const
     }
     for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
         auto& elementName = m_namedOutputs[i].elementName;
-        auto& internalTypeName = m_namedOutputs[i].internalTypeName;
         auto& path = m_entryPointItems.outputs[i].path;
-        stringBuilder.append(makeString(outputVariableName, '.', elementName, " = ", internalTypeName, '(', inputVariableName, mangledOutputPath(path), ");\n"));
+        stringBuilder.append(makeString(outputVariableName, '.', elementName, " = ", inputVariableName, mangledOutputPath(path), ";\n"));
     }
     return stringBuilder.toString();
 }
@@ -469,13 +424,8 @@ FragmentEntryPointScaffolding::FragmentEntryPointScaffolding(AST::FunctionDefini
 
     m_namedOutputs.reserveInitialCapacity(m_entryPointItems.outputs.size());
     for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
-        auto& outputItem = m_entryPointItems.outputs[i];
         NamedOutput namedOutput;
         namedOutput.elementName = m_typeNamer.generateNextStructureElementName();
-        if (WTF::holds_alternative<AST::BuiltInSemantic>(*outputItem.semantic))
-            namedOutput.internalTypeName = internalTypeForSemantic(WTF::get<AST::BuiltInSemantic>(*outputItem.semantic));
-        if (namedOutput.internalTypeName.isNull())
-            namedOutput.internalTypeName = m_typeNamer.mangledNameForType(*outputItem.unnamedType);
         m_namedOutputs.uncheckedAppend(WTFMove(namedOutput));
     }
 }
@@ -496,10 +446,10 @@ String FragmentEntryPointScaffolding::helperTypes()
     stringBuilder.append(makeString("struct ", m_returnStructName, " {\n"));
     for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
         auto& outputItem = m_entryPointItems.outputs[i];
-        auto& internalTypeName = m_namedOutputs[i].internalTypeName;
+        auto mangledTypeName = m_typeNamer.mangledNameForType(*outputItem.unnamedType);
         auto elementName = m_namedOutputs[i].elementName;
         auto attribute = attributeForSemantic(*outputItem.semantic);
-        stringBuilder.append(makeString("    ", internalTypeName, ' ', elementName, ' ', attribute, ";\n"));
+        stringBuilder.append(makeString("    ", mangledTypeName, ' ', elementName, ' ', attribute, ";\n"));
     }
     stringBuilder.append("};\n\n");
 
@@ -548,9 +498,8 @@ String FragmentEntryPointScaffolding::pack(const String& inputVariableName, cons
     }
     for (size_t i = 0; i < m_entryPointItems.outputs.size(); ++i) {
         auto& elementName = m_namedOutputs[i].elementName;
-        auto& internalTypeName = m_namedOutputs[i].internalTypeName;
         auto& path = m_entryPointItems.outputs[i].path;
-        stringBuilder.append(makeString(outputVariableName, '.', elementName, " = ", internalTypeName, '(', inputVariableName, mangledOutputPath(path), ");\n"));
+        stringBuilder.append(makeString(outputVariableName, '.', elementName, " = ", inputVariableName, mangledOutputPath(path), ";\n"));
     }
     return stringBuilder.toString();
 }
@@ -569,7 +518,7 @@ String ComputeEntryPointScaffolding::signature(String& functionName)
 {
     StringBuilder stringBuilder;
 
-    stringBuilder.append(makeString("kernel void ", functionName, '('));
+    stringBuilder.append(makeString("compute void ", functionName, '('));
     bool empty = true;
     if (auto resourceSignature = this->resourceSignature()) {
         empty = false;
index 0f5cea5..664f7fb 100644 (file)
@@ -129,7 +129,6 @@ private:
 
     struct NamedOutput {
         String elementName;
-        String internalTypeName;
     };
     Vector<NamedOutput> m_namedOutputs;
 };
@@ -158,7 +157,6 @@ private:
 
     struct NamedOutput {
         String elementName;
-        String internalTypeName;
     };
     Vector<NamedOutput> m_namedOutputs;
 };
diff --git a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLComputeDimensions.cpp b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLComputeDimensions.cpp
deleted file mode 100644 (file)
index c4b7d9c..0000000
+++ /dev/null
@@ -1,90 +0,0 @@
-/*
- * 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 "WHLSLComputeDimensions.h"
-
-#if ENABLE(WEBGPU)
-
-#include "WHLSLFunctionDeclaration.h"
-#include "WHLSLPrepare.h"
-#include "WHLSLProgram.h"
-#include <wtf/Optional.h>
-
-namespace WebCore {
-
-namespace WHLSL {
-
-class ComputeDimensionsVisitor : public Visitor {
-public:
-    ComputeDimensionsVisitor(AST::FunctionDefinition& entryPoint)
-        : m_entryPoint(entryPoint)
-    {
-    }
-
-    virtual ~ComputeDimensionsVisitor() = default;
-
-    Optional<ComputeDimensions> computeDimensions() const { return m_computeDimensions; }
-
-private:
-    void visit(AST::FunctionDeclaration& functionDeclaration) override
-    {
-        bool foundNumThreadsFunctionAttribute = false;
-        for (auto& functionAttribute : functionDeclaration.attributeBlock()) {
-            auto success = WTF::visit(WTF::makeVisitor([&](AST::NumThreadsFunctionAttribute& numThreadsFunctionAttribute) {
-                if (foundNumThreadsFunctionAttribute)
-                    return false;
-                foundNumThreadsFunctionAttribute = true;
-                if (&functionDeclaration == &m_entryPoint) {
-                    ASSERT(!m_computeDimensions);
-                    m_computeDimensions = {{ numThreadsFunctionAttribute.width(), numThreadsFunctionAttribute.height(), numThreadsFunctionAttribute.depth() }};
-                }
-                return true;
-            }), functionAttribute);
-            if (!success) {
-                setError();
-                return;
-            }
-        }
-    }
-
-    AST::FunctionDefinition& m_entryPoint;
-    Optional<ComputeDimensions> m_computeDimensions;
-};
-
-Optional<ComputeDimensions> computeDimensions(Program& program, AST::FunctionDefinition& entryPoint)
-{
-    ComputeDimensionsVisitor computeDimensions(entryPoint);
-    computeDimensions.Visitor::visit(program);
-    if (computeDimensions.error())
-        return WTF::nullopt;
-    return computeDimensions.computeDimensions();
-}
-
-} // namespace WHLSL
-
-} // namespace WebCore
-
-#endif // ENABLE(WEBGPU)
diff --git a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLComputeDimensions.h b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLComputeDimensions.h
deleted file mode 100644 (file)
index de45081..0000000
+++ /dev/null
@@ -1,44 +0,0 @@
-/*
- * 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 "WHLSLPrepare.h"
-
-namespace WebCore {
-
-namespace WHLSL {
-
-class Program;
-
-Optional<ComputeDimensions> computeDimensions(Program&, AST::FunctionDefinition&);
-
-}
-
-}
-
-#endif
index aa249c0..996dd84 100644 (file)
@@ -175,8 +175,7 @@ Optional<EntryPointItems> gatherEntryPointItems(const Intrinsics& intrinsics, AS
             return WTF::nullopt;
     }
     Gatherer outputGatherer(intrinsics, functionDefinition.semantic() ? &*functionDefinition.semantic() : nullptr);
-    if (*functionDefinition.entryPointType() != AST::EntryPointType::Compute)
-        outputGatherer.checkErrorAndVisit(functionDefinition.type());
+    outputGatherer.checkErrorAndVisit(functionDefinition.type());
     if (outputGatherer.error())
         return WTF::nullopt;
 
index d6b2e74..42848c7 100644 (file)
@@ -32,7 +32,6 @@
 #include "WHLSLAutoInitializeVariables.h"
 #include "WHLSLCheckDuplicateFunctions.h"
 #include "WHLSLChecker.h"
-#include "WHLSLComputeDimensions.h"
 #include "WHLSLFunctionStageChecker.h"
 #include "WHLSLHighZombieFinder.h"
 #include "WHLSLLiteralTypeChecker.h"
@@ -171,16 +170,12 @@ Optional<ComputePrepareResult> prepare(String& whlslSource, ComputePipelineDescr
     auto matchedSemantics = matchSemantics(*program, computePipelineDescriptor);
     if (!matchedSemantics)
         return WTF::nullopt;
-    auto computeDimensions = WHLSL::computeDimensions(*program, *matchedSemantics->shader);
-    if (!computeDimensions)
-        return WTF::nullopt;
 
     auto generatedCode = Metal::generateMetalCode(*program, WTFMove(*matchedSemantics), computePipelineDescriptor.layout);
 
     ComputePrepareResult result;
     result.metalSource = WTFMove(generatedCode.metalSource);
     result.mangledEntryPointName = WTFMove(generatedCode.mangledEntryPointName);
-    result.computeDimensions = WTFMove(*computeDimensions);
     return result;
 }
 
index eb2021b..5bd6df0 100644 (file)
@@ -42,16 +42,9 @@ struct RenderPrepareResult {
 };
 Optional<RenderPrepareResult> prepare(String& whlslSource, RenderPipelineDescriptor&);
 
-struct ComputeDimensions {
-    unsigned width;
-    unsigned height;
-    unsigned depth;
-};
-
 struct ComputePrepareResult {
     String metalSource;
     String mangledEntryPointName;
-    ComputeDimensions computeDimensions;
 };
 Optional<ComputePrepareResult> prepare(String& whlslSource, ComputePipelineDescriptor&);
 
index f80da99..a93e037 100644 (file)
@@ -774,8 +774,6 @@ void LeftValueSimplifier::finishVisiting(AST::PropertyAccessExpression& property
     ASSERT(propertyAccessExpression.base().typeAnnotation().leftAddressSpace());
     ASSERT(propertyAccessExpression.anderFunction());
 
-    Visitor::visit(propertyAccessExpression.base());
-
     Lexer::Token origin = propertyAccessExpression.origin();
     auto* anderFunction = propertyAccessExpression.anderFunction();
 
@@ -806,6 +804,7 @@ void LeftValueSimplifier::visit(AST::DotExpression& dotExpression)
 
 void LeftValueSimplifier::visit(AST::IndexExpression& indexExpression)
 {
+    Visitor::visit(indexExpression);
     PropertyResolver().Visitor::visit(indexExpression.indexExpression());
     finishVisiting(indexExpression);
 }
index 5170e1a..1c4527a 100644 (file)
@@ -432,178 +432,23 @@ native operator float(short);
 native operator float(int);
 native operator float(half);
 
-native float operator+(float, float);
-native float operator-(float, float);
-native int operator+(int, int);
-native uint operator+(uint, uint);
-native bool operator<(int, int);
-native bool operator<(uint, uint);
-native bool operator<(float, float);
-native float operator*(float, float);
-
-native bool operator.x(bool2);
-native bool operator.y(bool2);
-native bool operator.x(bool3);
-native bool operator.y(bool3);
-native bool operator.z(bool3);
-native bool operator.x(bool4);
-native bool operator.y(bool4);
-native bool operator.z(bool4);
-native bool operator.w(bool4);
-native bool2 operator.x=(bool2, bool);
-native bool2 operator.y=(bool2, bool);
-native bool3 operator.x=(bool3, bool);
-native bool3 operator.y=(bool3, bool);
-native bool3 operator.z=(bool3, bool);
-native bool4 operator.x=(bool4, bool);
-native bool4 operator.y=(bool4, bool);
-native bool4 operator.z=(bool4, bool);
-native bool4 operator.w=(bool4, bool);
-native uchar operator.x(uchar2);
-native uchar operator.y(uchar2);
-native uchar operator.x(uchar3);
-native uchar operator.y(uchar3);
-native uchar operator.z(uchar3);
-native uchar operator.x(uchar4);
-native uchar operator.y(uchar4);
-native uchar operator.z(uchar4);
-native uchar operator.w(uchar4);
-native uchar2 operator.x=(uchar2, uchar);
-native uchar2 operator.y=(uchar2, uchar);
-native uchar3 operator.x=(uchar3, uchar);
-native uchar3 operator.y=(uchar3, uchar);
-native uchar3 operator.z=(uchar3, uchar);
-native uchar4 operator.x=(uchar4, uchar);
-native uchar4 operator.y=(uchar4, uchar);
-native uchar4 operator.z=(uchar4, uchar);
-native uchar4 operator.w=(uchar4, uchar);
-native ushort operator.x(ushort2);
-native ushort operator.y(ushort2);
-native ushort operator.x(ushort3);
-native ushort operator.y(ushort3);
-native ushort operator.z(ushort3);
-native ushort operator.x(ushort4);
-native ushort operator.y(ushort4);
-native ushort operator.z(ushort4);
-native ushort operator.w(ushort4);
-native ushort2 operator.x=(ushort2, ushort);
-native ushort2 operator.y=(ushort2, ushort);
-native ushort3 operator.x=(ushort3, ushort);
-native ushort3 operator.y=(ushort3, ushort);
-native ushort3 operator.z=(ushort3, ushort);
-native ushort4 operator.x=(ushort4, ushort);
-native ushort4 operator.y=(ushort4, ushort);
-native ushort4 operator.z=(ushort4, ushort);
-native ushort4 operator.w=(ushort4, ushort);
-native uint operator.x(uint2);
-native uint operator.y(uint2);
-native uint operator.x(uint3);
-native uint operator.y(uint3);
-native uint operator.z(uint3);
-native uint operator.x(uint4);
-native uint operator.y(uint4);
-native uint operator.z(uint4);
-native uint operator.w(uint4);
-native uint2 operator.x=(uint2, uint);
-native uint2 operator.y=(uint2, uint);
-native uint3 operator.x=(uint3, uint);
-native uint3 operator.y=(uint3, uint);
-native uint3 operator.z=(uint3, uint);
-native uint4 operator.x=(uint4, uint);
-native uint4 operator.y=(uint4, uint);
-native uint4 operator.z=(uint4, uint);
-native uint4 operator.w=(uint4, uint);
-native char operator.x(char2);
-native char operator.y(char2);
-native char operator.x(char3);
-native char operator.y(char3);
-native char operator.z(char3);
-native char operator.x(char4);
-native char operator.y(char4);
-native char operator.z(char4);
-native char operator.w(char4);
-native char2 operator.x=(char2, char);
-native char2 operator.y=(char2, char);
-native char3 operator.x=(char3, char);
-native char3 operator.y=(char3, char);
-native char3 operator.z=(char3, char);
-native char4 operator.x=(char4, char);
-native char4 operator.y=(char4, char);
-native char4 operator.z=(char4, char);
-native char4 operator.w=(char4, char);
-native short operator.x(short2);
-native short operator.y(short2);
-native short operator.x(short3);
-native short operator.y(short3);
-native short operator.z(short3);
-native short operator.x(short4);
-native short operator.y(short4);
-native short operator.z(short4);
-native short operator.w(short4);
-native short2 operator.x=(short2, short);
-native short2 operator.y=(short2, short);
-native short3 operator.x=(short3, short);
-native short3 operator.y=(short3, short);
-native short3 operator.z=(short3, short);
-native short4 operator.x=(short4, short);
-native short4 operator.y=(short4, short);
-native short4 operator.z=(short4, short);
-native short4 operator.w=(short4, short);
-native int operator.x(int2);
-native int operator.y(int2);
-native int operator.x(int3);
-native int operator.y(int3);
-native int operator.z(int3);
-native int operator.x(int4);
-native int operator.y(int4);
-native int operator.z(int4);
-native int operator.w(int4);
-native int2 operator.x=(int2, int);
-native int2 operator.y=(int2, int);
-native int3 operator.x=(int3, int);
-native int3 operator.y=(int3, int);
-native int3 operator.z=(int3, int);
-native int4 operator.x=(int4, int);
-native int4 operator.y=(int4, int);
-native int4 operator.z=(int4, int);
-native int4 operator.w=(int4, int);
-native half operator.x(half2);
-native half operator.y(half2);
-native half operator.x(half3);
-native half operator.y(half3);
-native half operator.z(half3);
-native half operator.x(half4);
-native half operator.y(half4);
-native half operator.z(half4);
-native half operator.w(half4);
-native half2 operator.x=(half2, half);
-native half2 operator.y=(half2, half);
-native half3 operator.x=(half3, half);
-native half3 operator.y=(half3, half);
-native half3 operator.z=(half3, half);
-native half4 operator.x=(half4, half);
-native half4 operator.y=(half4, half);
-native half4 operator.z=(half4, half);
-native half4 operator.w=(half4, half);
-native float operator.x(float2);
-native float operator.y(float2);
-native float operator.x(float3);
-native float operator.y(float3);
-native float operator.z(float3);
 native float operator.x(float4);
 native float operator.y(float4);
 native float operator.z(float4);
 native float operator.w(float4);
-native float2 operator.x=(float2, float);
-native float2 operator.y=(float2, float);
-native float3 operator.x=(float3, float);
-native float3 operator.y=(float3, float);
-native float3 operator.z=(float3, float);
 native float4 operator.x=(float4, float);
 native float4 operator.y=(float4, float);
 native float4 operator.z=(float4, float);
 native float4 operator.w=(float4, float);
 
+native float operator+(float, float);
+native float operator-(float, float);
+native int operator+(int, int);
+native uint operator+(uint, uint);
+native bool operator<(int, int);
+native bool operator<(uint, uint);
+native bool operator<(float, float);
+
 native float ddx(float);
 native float ddy(float);
 native void AllMemoryBarrierWithGroupSync();
index 1b2cfd1..0d381d3 100644 (file)
@@ -306,7 +306,6 @@ Modules/websockets/WorkerThreadableWebSocketChannel.cpp
 
 Modules/webgpu/GPUCanvasContext.cpp
 Modules/webgpu/NavigatorGPU.cpp
-Modules/webgpu/WHLSL/WHLSLComputeDimensions.cpp
 Modules/webgpu/WHLSL/WHLSLASTDumper.cpp
 Modules/webgpu/WHLSL/WHLSLAutoInitializeVariables.cpp
 Modules/webgpu/WHLSL/WHLSLInferTypes.cpp
index 18bfd30..6048303 100644 (file)
@@ -326,7 +326,6 @@ platform/graphics/gpu/cocoa/GPUCommandBufferMetal.mm
 platform/graphics/gpu/cocoa/GPUComputePassEncoderMetal.mm
 platform/graphics/gpu/cocoa/GPUComputePipelineMetal.mm
 platform/graphics/gpu/cocoa/GPUDeviceMetal.mm
-platform/graphics/gpu/cocoa/GPUPipelineMetalConvertLayout.cpp
 platform/graphics/gpu/cocoa/GPUProgrammablePassEncoderMetal.mm
 platform/graphics/gpu/cocoa/GPUQueueMetal.mm
 platform/graphics/gpu/cocoa/GPURenderPassEncoderMetal.mm
index aacba4c..43007e6 100644 (file)
                1C840B9921EC400800D0500D /* WHLSLGatherEntryPointItems.cpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.cpp; path = WHLSLGatherEntryPointItems.cpp; sourceTree = "<group>"; };
                1C840B9A21EC400900D0500D /* WHLSLGatherEntryPointItems.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = WHLSLGatherEntryPointItems.h; sourceTree = "<group>"; };
                1C840B9B21EC400900D0500D /* WHLSLChecker.cpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.cpp; path = WHLSLChecker.cpp; sourceTree = "<group>"; };
-               1C86CA4B22AA19FF001BF961 /* WHLSLComputeDimensions.cpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.cpp; path = WHLSLComputeDimensions.cpp; sourceTree = "<group>"; };
-               1C86CA4C22AA19FF001BF961 /* WHLSLComputeDimensions.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = WHLSLComputeDimensions.h; sourceTree = "<group>"; };
-               1C86CA4E22AA23C9001BF961 /* GPUPipelineMetalConvertLayout.cpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.cpp; path = GPUPipelineMetalConvertLayout.cpp; sourceTree = "<group>"; };
-               1C86CA4F22AA23C9001BF961 /* GPUPipelineMetalConvertLayout.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = GPUPipelineMetalConvertLayout.h; sourceTree = "<group>"; };
                1C904DF90BA9D2C80081E9D0 /* Version.xcconfig */ = {isa = PBXFileReference; fileEncoding = 4; lastKnownFileType = text.xcconfig; path = Version.xcconfig; sourceTree = "<group>"; };
                1C9AE5CA21ED9DF50069D5F2 /* WHLSLHighZombieFinder.cpp */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.cpp.cpp; path = WHLSLHighZombieFinder.cpp; sourceTree = "<group>"; };
                1C9AE5CB21ED9DF50069D5F2 /* WHLSLHighZombieFinder.h */ = {isa = PBXFileReference; lastKnownFileType = sourcecode.c.h; path = WHLSLHighZombieFinder.h; sourceTree = "<group>"; };
                                C234A9AE21E92C1A003C984D /* WHLSLCheckDuplicateFunctions.h */,
                                1C840B9B21EC400900D0500D /* WHLSLChecker.cpp */,
                                1C840B9721EC400700D0500D /* WHLSLChecker.h */,
-                               1C86CA4B22AA19FF001BF961 /* WHLSLComputeDimensions.cpp */,
-                               1C86CA4C22AA19FF001BF961 /* WHLSLComputeDimensions.h */,
                                1CA0C2E421EED12A00A11860 /* WHLSLFunctionStageChecker.cpp */,
                                1CA0C2E521EED12A00A11860 /* WHLSLFunctionStageChecker.h */,
                                1C840B9921EC400800D0500D /* WHLSLGatherEntryPointItems.cpp */,
                                D08903402241CE4600F3F440 /* GPUComputePassEncoderMetal.mm */,
                                D089033B224179B500F3F440 /* GPUComputePipelineMetal.mm */,
                                D087CE3C21ACA94200BDE174 /* GPUDeviceMetal.mm */,
-                               1C86CA4E22AA23C9001BF961 /* GPUPipelineMetalConvertLayout.cpp */,
-                               1C86CA4F22AA23C9001BF961 /* GPUPipelineMetalConvertLayout.h */,
                                D087CE3B21ACA94200BDE174 /* GPUProgrammablePassEncoderMetal.mm */,
                                D087CE3921ACA94200BDE174 /* GPUQueueMetal.mm */,
                                D087CE3A21ACA94200BDE174 /* GPURenderPassEncoderMetal.mm */,
index 65367c9..0359f29 100644 (file)
@@ -27,7 +27,6 @@
 
 #if ENABLE(WEBGPU)
 
-#include "WHLSLPrepare.h"
 #include <wtf/RefCounted.h>
 #include <wtf/RefPtr.h>
 #include <wtf/RetainPtr.h>
@@ -49,13 +48,10 @@ public:
 
     const PlatformComputePipeline* platformComputePipeline() const { return m_platformComputePipeline.get(); }
 
-    WHLSL::ComputeDimensions computeDimensions() const { return m_computeDimensions; }
-
 private:
-    GPUComputePipeline(PlatformComputePipelineSmartPtr&&, WHLSL::ComputeDimensions);
+    GPUComputePipeline(PlatformComputePipelineSmartPtr&&);
 
     PlatformComputePipelineSmartPtr m_platformComputePipeline;
-    WHLSL::ComputeDimensions m_computeDimensions { 0, 0, 0 };
 };
 
 } // namespace WebCore
index a5b7ca9..713c1b7 100644 (file)
@@ -92,11 +92,16 @@ void GPUComputePassEncoder::dispatch(unsigned x, unsigned y, unsigned z)
         return;
     }
 
-    ASSERT(m_pipeline->platformComputePipeline());
+    auto pipelineState = m_pipeline->platformComputePipeline();
+    ASSERT(pipelineState);
 
     BEGIN_BLOCK_OBJC_EXCEPTIONS;
 
-    auto threadsPerThreadgroup = MTLSizeMake(m_pipeline->computeDimensions().width, m_pipeline->computeDimensions().height, m_pipeline->computeDimensions().depth);
+    auto w = pipelineState.threadExecutionWidth;
+    auto h = pipelineState.maxTotalThreadsPerThreadgroup / w;
+
+    // FIXME: This should be gleaned from the shader if not using MSL. For now, use the docs' example calculation.
+    auto threadsPerThreadgroup = MTLSizeMake(w, h, 1);
 
     auto threadgroupsPerGrid = MTLSizeMake(x, y, z);
 
index d6eced4..fbe43f4 100644 (file)
 
 #import "GPUComputePipelineDescriptor.h"
 #import "GPUDevice.h"
-#import "GPUPipelineMetalConvertLayout.h"
 #import "Logging.h"
-#import "WHLSLPrepare.h"
 #import <Metal/Metal.h>
 #import <wtf/BlockObjCExceptions.h>
 
-namespace WebCore {
-
-static bool trySetMetalFunctions(const char* const functionName, MTLLibrary *computeMetalLibrary, MTLComputePipelineDescriptor *mtlDescriptor, const String& computeEntryPointName)
-{
-#if LOG_DISABLED
-    UNUSED_PARAM(functionName);
-#endif
-
-    BEGIN_BLOCK_OBJC_EXCEPTIONS;
-
-    if (!computeMetalLibrary) {
-        LOG(WebGPU, "%s: MTLLibrary for compute stage does not exist!", functionName);
-        return false;
-    }
-
-    auto function = adoptNS([computeMetalLibrary newFunctionWithName:computeEntryPointName]);
-    if (!function) {
-        LOG(WebGPU, "%s: Cannot create compute MTLFunction \"%s\"!", functionName, computeEntryPointName.utf8().data());
-        return false;
-    }
-
-    [mtlDescriptor setComputeFunction:function.get()];
-    return true;
-
-    END_BLOCK_OBJC_EXCEPTIONS;
+OBJC_PROTOCOL(MTLFunction);
 
-    return false;
-}
+namespace WebCore {
 
-static Optional<WHLSL::ComputeDimensions> trySetFunctions(const char* const functionName, const GPUPipelineStageDescriptor& computeStage, const GPUDevice& device, MTLComputePipelineDescriptor* mtlDescriptor, Optional<WHLSL::ComputePipelineDescriptor>& whlslDescriptor)
+static RetainPtr<MTLFunction> tryCreateMtlComputeFunction(const GPUPipelineStageDescriptor& stage)
 {
-#if LOG_DISABLED
-    UNUSED_PARAM(functionName);
-#endif
-    RetainPtr<MTLLibrary> computeLibrary;
-    String computeEntryPoint;
-
-    WHLSL::ComputeDimensions computeDimensions { 1, 1, 1 };
-
-    if (whlslDescriptor) {
-        // WHLSL functions are compiled to MSL first.
-        String whlslSource = computeStage.module->whlslSource();
-        ASSERT(!whlslSource.isNull());
-
-        whlslDescriptor->entryPointName = computeStage.entryPoint;
-
-        auto whlslCompileResult = WHLSL::prepare(whlslSource, *whlslDescriptor);
-        if (!whlslCompileResult)
-            return WTF::nullopt;
-        computeDimensions = whlslCompileResult->computeDimensions;
-
-        NSError *error = nil;
-
-        BEGIN_BLOCK_OBJC_EXCEPTIONS;
-        computeLibrary = adoptNS([device.platformDevice() newLibraryWithSource:whlslCompileResult->metalSource options:nil error:&error]);
-        END_BLOCK_OBJC_EXCEPTIONS;
-
-        ASSERT(computeLibrary);
-        // FIXME: https://bugs.webkit.org/show_bug.cgi?id=195771 Once we zero-fill variables, there should be no warnings, so we should be able to ASSERT(!error) here.
-
-        computeEntryPoint = whlslCompileResult->mangledEntryPointName;
-    } else {
-        computeLibrary = computeStage.module->platformShaderModule();
-        computeEntryPoint = computeStage.entryPoint;
+    if (!stage.module->platformShaderModule() || stage.entryPoint.isNull()) {
+        LOG(WebGPU, "GPUComputePipeline::tryCreate(): Invalid GPUShaderModule!");
+        return nullptr;
     }
 
-    if (trySetMetalFunctions(functionName, computeLibrary.get(), mtlDescriptor, computeEntryPoint))
-        return computeDimensions;
-    return WTF::nullopt;
-}
-
-struct ConvertResult {
-    RetainPtr<MTLComputePipelineDescriptor> pipelineDescriptor;
-    WHLSL::ComputeDimensions computeDimensions;
-};
-static Optional<ConvertResult> convertComputePipelineDescriptor(const char* const functionName, const GPUComputePipelineDescriptor& descriptor, const GPUDevice& device)
-{
-    RetainPtr<MTLComputePipelineDescriptor> mtlDescriptor;
+    RetainPtr<MTLFunction> function;
 
     BEGIN_BLOCK_OBJC_EXCEPTIONS;
-
-    mtlDescriptor = adoptNS([MTLComputePipelineDescriptor new]);
-
+    function = adoptNS([stage.module->platformShaderModule() newFunctionWithName:stage.entryPoint]);
     END_BLOCK_OBJC_EXCEPTIONS;
 
-    if (!mtlDescriptor) {
-        LOG(WebGPU, "%s: Error creating MTLDescriptor!", functionName);
-        return WTF::nullopt;
-    }
-
-    const auto& computeStage = descriptor.computeStage;
-
-    bool isWhlsl = !computeStage.module->whlslSource().isNull();
-
-    Optional<WHLSL::ComputePipelineDescriptor> whlslDescriptor;
-    if (isWhlsl)
-        whlslDescriptor = WHLSL::ComputePipelineDescriptor();
-
-    if (descriptor.layout && whlslDescriptor) {
-        if (auto layout = convertLayout(*descriptor.layout))
-            whlslDescriptor->layout = WTFMove(*layout);
-        else {
-            LOG(WebGPU, "%s: Error converting GPUPipelineLayout!", functionName);
-            return WTF::nullopt;
-        }
-    }
-
-    if (auto computeDimensions = trySetFunctions(functionName, computeStage, device, mtlDescriptor.get(), whlslDescriptor))
-        return {{ mtlDescriptor, *computeDimensions }};
+    if (!function)
+        LOG(WebGPU, "GPUComputePipeline::tryCreate(): Cannot create compute MTLFunction \"%s\"!", stage.entryPoint.utf8().data());
 
-    return WTF::nullopt;
+    return function;
 }
 
-struct CreateResult {
-    RetainPtr<MTLComputePipelineState> pipelineState;
-    WHLSL::ComputeDimensions computeDimensions;
-};
-static Optional<CreateResult> tryCreateMTLComputePipelineState(const char* const functionName, const GPUDevice& device, const GPUComputePipelineDescriptor& descriptor)
+static RetainPtr<MTLComputePipelineState> tryCreateMTLComputePipelineState(const GPUDevice& device, const GPUComputePipelineDescriptor& descriptor)
 {
     if (!device.platformDevice()) {
         LOG(WebGPU, "GPUComputePipeline::tryCreate(): Invalid GPUDevice!");
-        return WTF::nullopt;
+        return nullptr;
     }
 
-    auto convertResult = convertComputePipelineDescriptor(functionName, descriptor, device);
-    if (!convertResult)
-        return WTF::nullopt;
-    ASSERT(convertResult->pipelineDescriptor);
-    auto mtlDescriptor = convertResult->pipelineDescriptor;
+    auto computeFunction = tryCreateMtlComputeFunction(descriptor.computeStage);
+    if (!computeFunction)
+        return nullptr;
 
-    RetainPtr<MTLComputePipelineState> pipeline;
+    RetainPtr<MTLComputePipelineState> pipelineState;
+    NSError *error = nil;
 
     BEGIN_BLOCK_OBJC_EXCEPTIONS;
+    pipelineState = adoptNS([device.platformDevice() newComputePipelineStateWithFunction:computeFunction.get() error:&error]);
+    END_BLOCK_OBJC_EXCEPTIONS;
 
-    NSError *error = nil;
-    pipeline = adoptNS([device.platformDevice() newComputePipelineStateWithDescriptor:mtlDescriptor.get() options:MTLPipelineOptionNone reflection:nil error:&error]);
-    if (!pipeline) {
+    if (!pipelineState)
         LOG(WebGPU, "GPUComputePipeline::tryCreate(): %s!", error ? error.localizedDescription.UTF8String : "Unable to create MTLComputePipelineState!");
-        return WTF::nullopt;
-    }
 
-    END_BLOCK_OBJC_EXCEPTIONS;
-
-    return {{ pipeline, convertResult->computeDimensions }};
+    return pipelineState;
 }
 
 RefPtr<GPUComputePipeline> GPUComputePipeline::tryCreate(const GPUDevice& device, const GPUComputePipelineDescriptor& descriptor)
 {
-    const char* const functionName = "GPURenderPipeline::create()";
-
-    auto createResult = tryCreateMTLComputePipelineState(functionName, device, descriptor);
-    if (!createResult)
+    auto mtlPipeline = tryCreateMTLComputePipelineState(device, descriptor);
+    if (!mtlPipeline)
         return nullptr;
 
-    return adoptRef(new GPUComputePipeline(WTFMove(createResult->pipelineState), createResult->computeDimensions));
+    return adoptRef(new GPUComputePipeline(WTFMove(mtlPipeline)));
 }
 
-GPUComputePipeline::GPUComputePipeline(RetainPtr<MTLComputePipelineState>&& pipeline, WHLSL::ComputeDimensions computeDimensions)
+GPUComputePipeline::GPUComputePipeline(RetainPtr<MTLComputePipelineState>&& pipeline)
     : m_platformComputePipeline(WTFMove(pipeline))
-    , m_computeDimensions(computeDimensions)
 {
 }
 
diff --git a/Source/WebCore/platform/graphics/gpu/cocoa/GPUPipelineMetalConvertLayout.cpp b/Source/WebCore/platform/graphics/gpu/cocoa/GPUPipelineMetalConvertLayout.cpp
deleted file mode 100644 (file)
index bd57688..0000000
+++ /dev/null
@@ -1,94 +0,0 @@
-/*
- * 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 "GPUPipelineMetalConvertLayout.h"
-
-#include "GPUPipelineLayout.h"
-
-#if ENABLE(WEBGPU)
-
-namespace WebCore {
-
-static OptionSet<WHLSL::ShaderStage> convertShaderStageFlags(GPUShaderStageFlags flags)
-{
-    OptionSet<WHLSL::ShaderStage> result;
-    if (flags & GPUShaderStageBit::Flags::Vertex)
-        result.add(WHLSL::ShaderStage::Vertex);
-    if (flags & GPUShaderStageBit::Flags::Fragment)
-        result.add(WHLSL::ShaderStage::Fragment);
-    if (flags & GPUShaderStageBit::Flags::Compute)
-        result.add(WHLSL::ShaderStage::Compute);
-    return result;
-}
-
-static Optional<WHLSL::Binding::BindingDetails> convertBindingType(GPUBindGroupLayout::InternalBindingDetails internalBindingDetails)
-{
-    return WTF::visit(WTF::makeVisitor([&](GPUBindGroupLayout::UniformBuffer uniformBuffer) -> Optional<WHLSL::Binding::BindingDetails> {
-        return { WHLSL::UniformBufferBinding { uniformBuffer.internalLengthName } };
-    }, [&](GPUBindGroupLayout::DynamicUniformBuffer) -> Optional<WHLSL::Binding::BindingDetails> {
-        return WTF::nullopt;
-    }, [&](GPUBindGroupLayout::Sampler) -> Optional<WHLSL::Binding::BindingDetails> {
-        return { WHLSL::SamplerBinding { } };
-    }, [&](GPUBindGroupLayout::SampledTexture) -> Optional<WHLSL::Binding::BindingDetails> {
-        return { WHLSL::TextureBinding { } };
-    }, [&](GPUBindGroupLayout::StorageBuffer storageBuffer) -> Optional<WHLSL::Binding::BindingDetails> {
-        return { WHLSL::StorageBufferBinding { storageBuffer.internalLengthName } };
-    }, [&](GPUBindGroupLayout::DynamicStorageBuffer) -> Optional<WHLSL::Binding::BindingDetails> {
-        return WTF::nullopt;
-    }), internalBindingDetails);
-}
-
-Optional<WHLSL::Layout> convertLayout(const GPUPipelineLayout& layout)
-{
-    WHLSL::Layout result;
-    if (layout.bindGroupLayouts().size() > std::numeric_limits<unsigned>::max())
-        return WTF::nullopt;
-    for (size_t i = 0; i < layout.bindGroupLayouts().size(); ++i) {
-        const auto& bindGroupLayout = layout.bindGroupLayouts()[i];
-        WHLSL::BindGroup bindGroup;
-        bindGroup.name = static_cast<unsigned>(i);
-        for (const auto& keyValuePair : bindGroupLayout->bindingsMap()) {
-            const auto& bindingDetails = keyValuePair.value;
-            WHLSL::Binding binding;
-            binding.visibility = convertShaderStageFlags(bindingDetails.externalBinding.visibility);
-            if (auto bindingType = convertBindingType(bindingDetails.internalBindingDetails))
-                binding.binding = *bindingType;
-            else
-                return WTF::nullopt;
-            if (bindingDetails.externalBinding.binding > std::numeric_limits<unsigned>::max())
-                return WTF::nullopt;
-            binding.externalName = bindingDetails.externalBinding.binding;
-            binding.internalName = bindingDetails.internalName;
-            bindGroup.bindings.append(WTFMove(binding));
-        }
-        result.append(WTFMove(bindGroup));
-    }
-    return result;
-}
-
-} // namespace WebCore
-
-#endif // ENABLE(WEBGPU)
diff --git a/Source/WebCore/platform/graphics/gpu/cocoa/GPUPipelineMetalConvertLayout.h b/Source/WebCore/platform/graphics/gpu/cocoa/GPUPipelineMetalConvertLayout.h
deleted file mode 100644 (file)
index 8e29d93..0000000
+++ /dev/null
@@ -1,41 +0,0 @@
-/*
- * 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 "WHLSLPipelineDescriptor.h"
-#include <wtf/Optional.h>
-
-namespace WebCore {
-
-class GPUPipelineLayout;
-
-Optional<WHLSL::Layout> convertLayout(const GPUPipelineLayout&);
-
-}
-
-#endif
index 1745df0..4a9eebb 100644 (file)
@@ -30,7 +30,6 @@
 
 #import "GPUDevice.h"
 #import "GPULimits.h"
-#import "GPUPipelineMetalConvertLayout.h"
 #import "GPUUtils.h"
 #import "Logging.h"
 #import "WHLSLPrepare.h"
@@ -99,6 +98,35 @@ static WHLSL::VertexFormat convertVertexFormat(GPUVertexFormat vertexFormat)
     }
 }
 
+static OptionSet<WHLSL::ShaderStage> convertShaderStageFlags(GPUShaderStageFlags flags)
+{
+    OptionSet<WHLSL::ShaderStage> result;
+    if (flags & GPUShaderStageBit::Flags::Vertex)
+        result.add(WHLSL::ShaderStage::Vertex);
+    if (flags & GPUShaderStageBit::Flags::Fragment)
+        result.add(WHLSL::ShaderStage::Fragment);
+    if (flags & GPUShaderStageBit::Flags::Compute)
+        result.add(WHLSL::ShaderStage::Compute);
+    return result;
+}
+
+static Optional<WHLSL::Binding::BindingDetails> convertBindingType(GPUBindGroupLayout::InternalBindingDetails internalBindingDetails)
+{
+    return WTF::visit(WTF::makeVisitor([&](GPUBindGroupLayout::UniformBuffer uniformBuffer) -> Optional<WHLSL::Binding::BindingDetails> {
+        return { WHLSL::UniformBufferBinding { uniformBuffer.internalLengthName } };
+    }, [&](GPUBindGroupLayout::DynamicUniformBuffer) -> Optional<WHLSL::Binding::BindingDetails> {
+        return WTF::nullopt;
+    }, [&](GPUBindGroupLayout::Sampler) -> Optional<WHLSL::Binding::BindingDetails> {
+        return { WHLSL::SamplerBinding { } };
+    }, [&](GPUBindGroupLayout::SampledTexture) -> Optional<WHLSL::Binding::BindingDetails> {
+        return { WHLSL::TextureBinding { } };
+    }, [&](GPUBindGroupLayout::StorageBuffer storageBuffer) -> Optional<WHLSL::Binding::BindingDetails> {
+        return { WHLSL::StorageBufferBinding { storageBuffer.internalLengthName } };
+    }, [&](GPUBindGroupLayout::DynamicStorageBuffer) -> Optional<WHLSL::Binding::BindingDetails> {
+        return WTF::nullopt;
+    }), internalBindingDetails);
+}
+
 static Optional<WHLSL::TextureFormat> convertTextureFormat(GPUTextureFormat format)
 {
     switch (format) {
@@ -340,6 +368,34 @@ static bool trySetColorStates(const char* const functionName, const Vector<GPUCo
     return true;
 }
 
+static Optional<WHLSL::Layout> convertLayout(const GPUPipelineLayout& layout)
+{
+    WHLSL::Layout result;
+    if (layout.bindGroupLayouts().size() > std::numeric_limits<unsigned>::max())
+        return WTF::nullopt;
+    for (size_t i = 0; i < layout.bindGroupLayouts().size(); ++i) {
+        const auto& bindGroupLayout = layout.bindGroupLayouts()[i];
+        WHLSL::BindGroup bindGroup;
+        bindGroup.name = static_cast<unsigned>(i);
+        for (const auto& keyValuePair : bindGroupLayout->bindingsMap()) {
+            const auto& bindingDetails = keyValuePair.value;
+            WHLSL::Binding binding;
+            binding.visibility = convertShaderStageFlags(bindingDetails.externalBinding.visibility);
+            if (auto bindingType = convertBindingType(bindingDetails.internalBindingDetails))
+                binding.binding = *bindingType;
+            else
+                return WTF::nullopt;
+            if (bindingDetails.externalBinding.binding > std::numeric_limits<unsigned>::max())
+                return WTF::nullopt;
+            binding.externalName = bindingDetails.externalBinding.binding;
+            binding.internalName = bindingDetails.internalName;
+            bindGroup.bindings.append(WTFMove(binding));
+        }
+        result.append(WTFMove(bindGroup));
+    }
+    return result;
+}
+
 static bool trySetMetalFunctions(const char* const functionName, MTLLibrary *vertexMetalLibrary, MTLLibrary *fragmentMetalLibrary, MTLRenderPipelineDescriptor *mtlDescriptor, const String& vertexEntryPointName, const String& fragmentEntryPointName)
 {
 #if LOG_DISABLED
@@ -410,6 +466,8 @@ static bool trySetFunctions(const char* const functionName, const GPUPipelineSta
         if (!whlslCompileResult)
             return false;
 
+        WTFLogAlways("Metal Source: %s", whlslCompileResult->metalSource.utf8().data());
+
         NSError *error = nil;
 
         BEGIN_BLOCK_OBJC_EXCEPTIONS;
@@ -484,11 +542,6 @@ static RetainPtr<MTLRenderPipelineDescriptor> convertRenderPipelineDescriptor(co
 
 static RetainPtr<MTLRenderPipelineState> tryCreateMtlRenderPipelineState(const char* const functionName, const GPURenderPipelineDescriptor& descriptor, const GPUDevice& device)
 {
-    if (!device.platformDevice()) {
-        LOG(WebGPU, "GPUComputePipeline::tryCreate(): Invalid GPUDevice!");
-        return nullptr;
-    }
-
     auto mtlDescriptor = convertRenderPipelineDescriptor(functionName, descriptor, device);
     if (!mtlDescriptor)
         return nullptr;
@@ -497,7 +550,7 @@ static RetainPtr<MTLRenderPipelineState> tryCreateMtlRenderPipelineState(const c
 
     BEGIN_BLOCK_OBJC_EXCEPTIONS;
 
-    NSError *error = nil;
+    NSError *error = [NSError errorWithDomain:@"com.apple.WebKit.GPU" code:1 userInfo:nil];
     pipeline = adoptNS([device.platformDevice() newRenderPipelineStateWithDescriptor:mtlDescriptor.get() error:&error]);
     if (!pipeline)
         LOG(WebGPU, "%s: %s!", functionName, error.localizedDescription.UTF8String);