Diff
Modified: trunk/LayoutTests/ChangeLog (246408 => 246409)
--- trunk/LayoutTests/ChangeLog 2019-06-13 18:47:22 UTC (rev 246408)
+++ trunk/LayoutTests/ChangeLog 2019-06-13 19:06:31 UTC (rev 246409)
@@ -1,3 +1,20 @@
+2019-06-13 Commit Queue <[email protected]>
+
+ 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 <[email protected]>
twitch.tv: embedded video hovers down the screen when scrolling on iPad
Added: trunk/LayoutTests/webgpu/compute-squares-expected.txt (0 => 246409)
--- trunk/LayoutTests/webgpu/compute-squares-expected.txt (rev 0)
+++ trunk/LayoutTests/webgpu/compute-squares-expected.txt 2019-06-13 19:06:31 UTC (rev 246409)
@@ -0,0 +1 @@
+
Added: trunk/LayoutTests/webgpu/compute-squares.html (0 => 246409)
--- trunk/LayoutTests/webgpu/compute-squares.html (rev 0)
+++ trunk/LayoutTests/webgpu/compute-squares.html 2019-06-13 19:06:31 UTC (rev 246409)
@@ -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=""
+<script src=""
+<script src=""
+<script>
+if (window.testRunner)
+ testRunner.waitUntilDone();
+
+const data = "" 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
Deleted: trunk/LayoutTests/webgpu/whlsl-compute-expected.txt (246408 => 246409)
--- trunk/LayoutTests/webgpu/whlsl-compute-expected.txt 2019-06-13 18:47:22 UTC (rev 246408)
+++ trunk/LayoutTests/webgpu/whlsl-compute-expected.txt 2019-06-13 19:06:31 UTC (rev 246409)
@@ -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
-
Deleted: trunk/LayoutTests/webgpu/whlsl-compute.html (246408 => 246409)
--- trunk/LayoutTests/webgpu/whlsl-compute.html 2019-06-13 18:47:22 UTC (rev 246408)
+++ trunk/LayoutTests/webgpu/whlsl-compute.html 2019-06-13 19:06:31 UTC (rev 246409)
@@ -1,88 +0,0 @@
-<!DOCTYPE html>
-<html>
-<head>
-<script src=""
-</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=""
-</body>
-</html>
Modified: trunk/Source/WebCore/ChangeLog (246408 => 246409)
--- trunk/Source/WebCore/ChangeLog 2019-06-13 18:47:22 UTC (rev 246408)
+++ trunk/Source/WebCore/ChangeLog 2019-06-13 19:06:31 UTC (rev 246409)
@@ -1,3 +1,20 @@
+2019-06-13 Commit Queue <[email protected]>
+
+ 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 <[email protected]>
twitch.tv: embedded video hovers down the screen when scrolling on iPad
Modified: trunk/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp (246408 => 246409)
--- trunk/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp 2019-06-13 18:47:22 UTC (rev 246408)
+++ trunk/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp 2019-06-13 19:06:31 UTC (rev 246409)
@@ -178,41 +178,6 @@
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 @@
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 @@
}
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 @@
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 @@
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 @@
}
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 @@
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 @@
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 @@
}
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 @@
{
StringBuilder stringBuilder;
- stringBuilder.append(makeString("kernel void ", functionName, '('));
+ stringBuilder.append(makeString("compute void ", functionName, '('));
bool empty = true;
if (auto resourceSignature = this->resourceSignature()) {
empty = false;
Modified: trunk/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.h (246408 => 246409)
--- trunk/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.h 2019-06-13 18:47:22 UTC (rev 246408)
+++ trunk/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.h 2019-06-13 19:06:31 UTC (rev 246409)
@@ -129,7 +129,6 @@
struct NamedOutput {
String elementName;
- String internalTypeName;
};
Vector<NamedOutput> m_namedOutputs;
};
@@ -158,7 +157,6 @@
struct NamedOutput {
String elementName;
- String internalTypeName;
};
Vector<NamedOutput> m_namedOutputs;
};
Deleted: trunk/Source/WebCore/Modules/webgpu/WHLSL/WHLSLComputeDimensions.cpp (246408 => 246409)
--- trunk/Source/WebCore/Modules/webgpu/WHLSL/WHLSLComputeDimensions.cpp 2019-06-13 18:47:22 UTC (rev 246408)
+++ trunk/Source/WebCore/Modules/webgpu/WHLSL/WHLSLComputeDimensions.cpp 2019-06-13 19:06:31 UTC (rev 246409)
@@ -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)
Deleted: trunk/Source/WebCore/Modules/webgpu/WHLSL/WHLSLComputeDimensions.h (246408 => 246409)
--- trunk/Source/WebCore/Modules/webgpu/WHLSL/WHLSLComputeDimensions.h 2019-06-13 18:47:22 UTC (rev 246408)
+++ trunk/Source/WebCore/Modules/webgpu/WHLSL/WHLSLComputeDimensions.h 2019-06-13 19:06:31 UTC (rev 246409)
@@ -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
Modified: trunk/Source/WebCore/Modules/webgpu/WHLSL/WHLSLGatherEntryPointItems.cpp (246408 => 246409)
--- trunk/Source/WebCore/Modules/webgpu/WHLSL/WHLSLGatherEntryPointItems.cpp 2019-06-13 18:47:22 UTC (rev 246408)
+++ trunk/Source/WebCore/Modules/webgpu/WHLSL/WHLSLGatherEntryPointItems.cpp 2019-06-13 19:06:31 UTC (rev 246409)
@@ -175,8 +175,7 @@
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;
Modified: trunk/Source/WebCore/Modules/webgpu/WHLSL/WHLSLPrepare.cpp (246408 => 246409)
--- trunk/Source/WebCore/Modules/webgpu/WHLSL/WHLSLPrepare.cpp 2019-06-13 18:47:22 UTC (rev 246408)
+++ trunk/Source/WebCore/Modules/webgpu/WHLSL/WHLSLPrepare.cpp 2019-06-13 19:06:31 UTC (rev 246409)
@@ -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,9 +170,6 @@
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);
@@ -180,7 +176,6 @@
ComputePrepareResult result;
result.metalSource = WTFMove(generatedCode.metalSource);
result.mangledEntryPointName = WTFMove(generatedCode.mangledEntryPointName);
- result.computeDimensions = WTFMove(*computeDimensions);
return result;
}
Modified: trunk/Source/WebCore/Modules/webgpu/WHLSL/WHLSLPrepare.h (246408 => 246409)
--- trunk/Source/WebCore/Modules/webgpu/WHLSL/WHLSLPrepare.h 2019-06-13 18:47:22 UTC (rev 246408)
+++ trunk/Source/WebCore/Modules/webgpu/WHLSL/WHLSLPrepare.h 2019-06-13 19:06:31 UTC (rev 246409)
@@ -42,16 +42,9 @@
};
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&);
Modified: trunk/Source/WebCore/Modules/webgpu/WHLSL/WHLSLPropertyResolver.cpp (246408 => 246409)
--- trunk/Source/WebCore/Modules/webgpu/WHLSL/WHLSLPropertyResolver.cpp 2019-06-13 18:47:22 UTC (rev 246408)
+++ trunk/Source/WebCore/Modules/webgpu/WHLSL/WHLSLPropertyResolver.cpp 2019-06-13 19:06:31 UTC (rev 246409)
@@ -774,8 +774,6 @@
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::IndexExpression& indexExpression)
{
+ Visitor::visit(indexExpression);
PropertyResolver().Visitor::visit(indexExpression.indexExpression());
finishVisiting(indexExpression);
}
Modified: trunk/Source/WebCore/Modules/webgpu/WHLSL/WHLSLStandardLibrary.txt (246408 => 246409)
--- trunk/Source/WebCore/Modules/webgpu/WHLSL/WHLSLStandardLibrary.txt 2019-06-13 18:47:22 UTC (rev 246408)
+++ trunk/Source/WebCore/Modules/webgpu/WHLSL/WHLSLStandardLibrary.txt 2019-06-13 19:06:31 UTC (rev 246409)
@@ -432,178 +432,23 @@
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();
Modified: trunk/Source/WebCore/Sources.txt (246408 => 246409)
--- trunk/Source/WebCore/Sources.txt 2019-06-13 18:47:22 UTC (rev 246408)
+++ trunk/Source/WebCore/Sources.txt 2019-06-13 19:06:31 UTC (rev 246409)
@@ -306,7 +306,6 @@
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
Modified: trunk/Source/WebCore/SourcesCocoa.txt (246408 => 246409)
--- trunk/Source/WebCore/SourcesCocoa.txt 2019-06-13 18:47:22 UTC (rev 246408)
+++ trunk/Source/WebCore/SourcesCocoa.txt 2019-06-13 19:06:31 UTC (rev 246409)
@@ -326,7 +326,6 @@
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
Modified: trunk/Source/WebCore/WebCore.xcodeproj/project.pbxproj (246408 => 246409)
--- trunk/Source/WebCore/WebCore.xcodeproj/project.pbxproj 2019-06-13 18:47:22 UTC (rev 246408)
+++ trunk/Source/WebCore/WebCore.xcodeproj/project.pbxproj 2019-06-13 19:06:31 UTC (rev 246409)
@@ -6394,10 +6394,6 @@
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>"; };
@@ -25452,8 +25448,6 @@
C234A9AE21E92C1A003C984D /* WHLSLCheckDuplicateFunctions.h */,
1C840B9B21EC400900D0500D /* WHLSLChecker.cpp */,
1C840B9721EC400700D0500D /* WHLSLChecker.h */,
- 1C86CA4B22AA19FF001BF961 /* WHLSLComputeDimensions.cpp */,
- 1C86CA4C22AA19FF001BF961 /* WHLSLComputeDimensions.h */,
1CA0C2E421EED12A00A11860 /* WHLSLFunctionStageChecker.cpp */,
1CA0C2E521EED12A00A11860 /* WHLSLFunctionStageChecker.h */,
1C840B9921EC400800D0500D /* WHLSLGatherEntryPointItems.cpp */,
@@ -25993,8 +25987,6 @@
D08903402241CE4600F3F440 /* GPUComputePassEncoderMetal.mm */,
D089033B224179B500F3F440 /* GPUComputePipelineMetal.mm */,
D087CE3C21ACA94200BDE174 /* GPUDeviceMetal.mm */,
- 1C86CA4E22AA23C9001BF961 /* GPUPipelineMetalConvertLayout.cpp */,
- 1C86CA4F22AA23C9001BF961 /* GPUPipelineMetalConvertLayout.h */,
D087CE3B21ACA94200BDE174 /* GPUProgrammablePassEncoderMetal.mm */,
D087CE3921ACA94200BDE174 /* GPUQueueMetal.mm */,
D087CE3A21ACA94200BDE174 /* GPURenderPassEncoderMetal.mm */,
Modified: trunk/Source/WebCore/platform/graphics/gpu/GPUComputePipeline.h (246408 => 246409)
--- trunk/Source/WebCore/platform/graphics/gpu/GPUComputePipeline.h 2019-06-13 18:47:22 UTC (rev 246408)
+++ trunk/Source/WebCore/platform/graphics/gpu/GPUComputePipeline.h 2019-06-13 19:06:31 UTC (rev 246409)
@@ -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 @@
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
Modified: trunk/Source/WebCore/platform/graphics/gpu/cocoa/GPUComputePassEncoderMetal.mm (246408 => 246409)
--- trunk/Source/WebCore/platform/graphics/gpu/cocoa/GPUComputePassEncoderMetal.mm 2019-06-13 18:47:22 UTC (rev 246408)
+++ trunk/Source/WebCore/platform/graphics/gpu/cocoa/GPUComputePassEncoderMetal.mm 2019-06-13 19:06:31 UTC (rev 246409)
@@ -92,12 +92,17 @@
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);
[m_platformComputePassEncoder dispatchThreadgroups:threadgroupsPerGrid threadsPerThreadgroup:threadsPerThreadgroup];
Modified: trunk/Source/WebCore/platform/graphics/gpu/cocoa/GPUComputePipelineMetal.mm (246408 => 246409)
--- trunk/Source/WebCore/platform/graphics/gpu/cocoa/GPUComputePipelineMetal.mm 2019-06-13 18:47:22 UTC (rev 246408)
+++ trunk/Source/WebCore/platform/graphics/gpu/cocoa/GPUComputePipelineMetal.mm 2019-06-13 19:06:31 UTC (rev 246409)
@@ -30,172 +30,68 @@
#import "GPUComputePipelineDescriptor.h"
#import "GPUDevice.h"
-#import "GPUPipelineMetalConvertLayout.h"
#import "Logging.h"
-#import "WHLSLPrepare.h"
#import <Metal/Metal.h>
#import <wtf/BlockObjCExceptions.h>
+OBJC_PROTOCOL(MTLFunction);
+
namespace WebCore {
-static bool trySetMetalFunctions(const char* const functionName, MTLLibrary *computeMetalLibrary, MTLComputePipelineDescriptor *mtlDescriptor, const String& computeEntryPointName)
+static RetainPtr<MTLFunction> tryCreateMtlComputeFunction(const GPUPipelineStageDescriptor& stage)
{
-#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;
+ if (!stage.module->platformShaderModule() || stage.entryPoint.isNull()) {
+ LOG(WebGPU, "GPUComputePipeline::tryCreate(): Invalid GPUShaderModule!");
+ return nullptr;
}
- auto function = adoptNS([computeMetalLibrary newFunctionWithName:computeEntryPointName]);
- if (!function) {
- LOG(WebGPU, "%s: Cannot create compute MTLFunction \"%s\"!", functionName, computeEntryPointName.utf8().data());
- return false;
- }
+ RetainPtr<MTLFunction> function;
- [mtlDescriptor setComputeFunction:function.get()];
- return true;
-
- END_BLOCK_OBJC_EXCEPTIONS;
-
- return false;
-}
-
-static Optional<WHLSL::ComputeDimensions> trySetFunctions(const char* const functionName, const GPUPipelineStageDescriptor& computeStage, const GPUDevice& device, MTLComputePipelineDescriptor* mtlDescriptor, Optional<WHLSL::ComputePipelineDescriptor>& whlslDescriptor)
-{
-#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 (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;
-
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;
- }
+ if (!function)
+ LOG(WebGPU, "GPUComputePipeline::tryCreate(): Cannot create compute MTLFunction \"%s\"!", stage.entryPoint.utf8().data());
- 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 }};
-
- 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)
{
}
Deleted: trunk/Source/WebCore/platform/graphics/gpu/cocoa/GPUPipelineMetalConvertLayout.cpp (246408 => 246409)
--- trunk/Source/WebCore/platform/graphics/gpu/cocoa/GPUPipelineMetalConvertLayout.cpp 2019-06-13 18:47:22 UTC (rev 246408)
+++ trunk/Source/WebCore/platform/graphics/gpu/cocoa/GPUPipelineMetalConvertLayout.cpp 2019-06-13 19:06:31 UTC (rev 246409)
@@ -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)
Deleted: trunk/Source/WebCore/platform/graphics/gpu/cocoa/GPUPipelineMetalConvertLayout.h (246408 => 246409)
--- trunk/Source/WebCore/platform/graphics/gpu/cocoa/GPUPipelineMetalConvertLayout.h 2019-06-13 18:47:22 UTC (rev 246408)
+++ trunk/Source/WebCore/platform/graphics/gpu/cocoa/GPUPipelineMetalConvertLayout.h 2019-06-13 19:06:31 UTC (rev 246409)
@@ -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
Modified: trunk/Source/WebCore/platform/graphics/gpu/cocoa/GPURenderPipelineMetal.mm (246408 => 246409)
--- trunk/Source/WebCore/platform/graphics/gpu/cocoa/GPURenderPipelineMetal.mm 2019-06-13 18:47:22 UTC (rev 246408)
+++ trunk/Source/WebCore/platform/graphics/gpu/cocoa/GPURenderPipelineMetal.mm 2019-06-13 19:06:31 UTC (rev 246409)
@@ -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 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 @@
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 @@
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<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 @@
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);
