WebKit Bugzilla
Attachment 372073 Details for
Bug 198837
: REGRESSION(r246396): Breaks internal builds (Requested by ShawnRoberts on #webkit).
Home
|
New
|
Browse
|
Search
|
[?]
|
Reports
|
Requests
|
Help
|
New Account
|
Log In
Remember
[x]
|
Forgot Password
Login:
[x]
[patch]
ROLLOUT of r246396
bug-198837-20190613120455.patch (text/plain), 62.86 KB, created by
WebKit Commit Bot
on 2019-06-13 12:04:56 PDT
(
hide
)
Description:
ROLLOUT of r246396
Filename:
MIME Type:
Creator:
WebKit Commit Bot
Created:
2019-06-13 12:04:56 PDT
Size:
62.86 KB
patch
obsolete
>Subversion Revision: 246408 >diff --git a/Source/WebCore/ChangeLog b/Source/WebCore/ChangeLog >index e7dc3d9926b9f9889ad98bf9ba459a783e8c04d8..6ef0baa45f83432dbb0057d4a8e2b732eedfd607 100644 >--- a/Source/WebCore/ChangeLog >+++ b/Source/WebCore/ChangeLog >@@ -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/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp >index 3baa37fa54c755fd134ebc232457f7d2247449db..da8c057857e28d2e1fbfafcdaba512e9e2003f30 100644 >--- a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp >+++ b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp >@@ -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; >diff --git a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.h b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.h >index 0f5cea55e6b59cfe85859bc2161ba9ab518b46a2..664f7fb4a237619a3c2e50f172edfa0d4a5ee3d2 100644 >--- a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.h >+++ b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.h >@@ -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 >index c4b7d9c9500c25c49d14454cdd0554824b0403e6..0000000000000000000000000000000000000000 >--- a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLComputeDimensions.cpp >+++ /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 >index de450813a58c9957aafad4f66e23dd31e11b9873..0000000000000000000000000000000000000000 >--- a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLComputeDimensions.h >+++ /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 >diff --git a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLGatherEntryPointItems.cpp b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLGatherEntryPointItems.cpp >index aa249c02b72686151b7cd9856630d1c31ffce790..996dd84072d529fa08980d2c1a098ba165bfe54b 100644 >--- a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLGatherEntryPointItems.cpp >+++ b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLGatherEntryPointItems.cpp >@@ -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; > >diff --git a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLPrepare.cpp b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLPrepare.cpp >index d6b2e74d0b264e6c9fef1ad8cd1050ebdce0958e..42848c7101d66ce2a88f112b1773f35669447432 100644 >--- a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLPrepare.cpp >+++ b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLPrepare.cpp >@@ -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; > } > >diff --git a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLPrepare.h b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLPrepare.h >index eb2021b7ca2ec58fbd4a05494c54708310e50ed6..5bd6df06dc691a72f1ce24dfc55576ec7f5284ab 100644 >--- a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLPrepare.h >+++ b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLPrepare.h >@@ -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&); > >diff --git a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLPropertyResolver.cpp b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLPropertyResolver.cpp >index f80da99758d849a2ec189634f3f3244a4151956d..a93e0372566c83052294d0261a166fa6151ab658 100644 >--- a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLPropertyResolver.cpp >+++ b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLPropertyResolver.cpp >@@ -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); > } >diff --git a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLStandardLibrary.txt b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLStandardLibrary.txt >index 5170e1aecd5bbef45612236a76334b7dd3dbe5b3..1c4527a6bf146c4df57baa92e5e33f25e397bdc3 100644 >--- a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLStandardLibrary.txt >+++ b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLStandardLibrary.txt >@@ -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(); >diff --git a/Source/WebCore/Sources.txt b/Source/WebCore/Sources.txt >index 1b2cfd19487f8d62d530d75b54055d57bcf504af..0d381d38c1d6c9348cc36ae005828bcd6f40a993 100644 >--- a/Source/WebCore/Sources.txt >+++ b/Source/WebCore/Sources.txt >@@ -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 >diff --git a/Source/WebCore/SourcesCocoa.txt b/Source/WebCore/SourcesCocoa.txt >index 18bfd301d965632b7b1029e08ba361b11e649368..60483037d205fa32d82810e8924eeb7587fdadd7 100644 >--- a/Source/WebCore/SourcesCocoa.txt >+++ b/Source/WebCore/SourcesCocoa.txt >@@ -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 >diff --git a/Source/WebCore/WebCore.xcodeproj/project.pbxproj b/Source/WebCore/WebCore.xcodeproj/project.pbxproj >index aacba4c14b071add6cd196b1ea7925fe76e5fe72..43007e6b17560d502ab5edffe7d6fb6e9b138e81 100644 >--- a/Source/WebCore/WebCore.xcodeproj/project.pbxproj >+++ b/Source/WebCore/WebCore.xcodeproj/project.pbxproj >@@ -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 */, >diff --git a/Source/WebCore/platform/graphics/gpu/GPUComputePipeline.h b/Source/WebCore/platform/graphics/gpu/GPUComputePipeline.h >index 65367c9d9343a5c4691878a194c32d84087b1013..0359f29bed0601e778eef4fde2853daf0745f06c 100644 >--- a/Source/WebCore/platform/graphics/gpu/GPUComputePipeline.h >+++ b/Source/WebCore/platform/graphics/gpu/GPUComputePipeline.h >@@ -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 >diff --git a/Source/WebCore/platform/graphics/gpu/cocoa/GPUComputePassEncoderMetal.mm b/Source/WebCore/platform/graphics/gpu/cocoa/GPUComputePassEncoderMetal.mm >index a5b7ca963395bdb76694c94bf18fb0e148071fdc..713c1b791075a086b90b31309914a405d6afe682 100644 >--- a/Source/WebCore/platform/graphics/gpu/cocoa/GPUComputePassEncoderMetal.mm >+++ b/Source/WebCore/platform/graphics/gpu/cocoa/GPUComputePassEncoderMetal.mm >@@ -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); > >diff --git a/Source/WebCore/platform/graphics/gpu/cocoa/GPUComputePipelineMetal.mm b/Source/WebCore/platform/graphics/gpu/cocoa/GPUComputePipelineMetal.mm >index d6eced44f4b5a9130c56552718d4de57d66d59d7..fbe43f441c83a8445c0bc01bd537b081501ac79d 100644 >--- a/Source/WebCore/platform/graphics/gpu/cocoa/GPUComputePipelineMetal.mm >+++ b/Source/WebCore/platform/graphics/gpu/cocoa/GPUComputePipelineMetal.mm >@@ -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> > >-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 >index bd57688f5be187c664dd138fe24652fd110743ff..0000000000000000000000000000000000000000 >--- a/Source/WebCore/platform/graphics/gpu/cocoa/GPUPipelineMetalConvertLayout.cpp >+++ /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 >index 8e29d93af025ea2e1a611ccaf392e00cc16bb620..0000000000000000000000000000000000000000 >--- a/Source/WebCore/platform/graphics/gpu/cocoa/GPUPipelineMetalConvertLayout.h >+++ /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 >diff --git a/Source/WebCore/platform/graphics/gpu/cocoa/GPURenderPipelineMetal.mm b/Source/WebCore/platform/graphics/gpu/cocoa/GPURenderPipelineMetal.mm >index 1745df017b67dd4ec8020c40cead40ec3d6765c3..4a9eebbe1302660f0e9b644b38e5ac83b848a7e9 100644 >--- a/Source/WebCore/platform/graphics/gpu/cocoa/GPURenderPipelineMetal.mm >+++ b/Source/WebCore/platform/graphics/gpu/cocoa/GPURenderPipelineMetal.mm >@@ -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); >diff --git a/LayoutTests/ChangeLog b/LayoutTests/ChangeLog >index 5d4d28b0362fce2f640aca32479f59f21b18b37f..9ba39642e682c543c6443460039c7881102421f0 100644 >--- a/LayoutTests/ChangeLog >+++ b/LayoutTests/ChangeLog >@@ -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 >index 0000000000000000000000000000000000000000..8b137891791fe96927ad78e64b0aad7bded08bdc >--- /dev/null >+++ b/LayoutTests/webgpu/compute-squares-expected.txt >@@ -0,0 +1 @@ >+ >diff --git a/LayoutTests/webgpu/compute-squares.html b/LayoutTests/webgpu/compute-squares.html >new file mode 100644 >index 0000000000000000000000000000000000000000..39bffc08d2ab0961c4eaa65e69224fb6bc5ca320 >--- /dev/null >+++ b/LayoutTests/webgpu/compute-squares.html >@@ -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 >index 6e1668a103beb92f76ea7e8d7aa7ffb50e89e241..0000000000000000000000000000000000000000 >--- a/LayoutTests/webgpu/whlsl-compute-expected.txt >+++ /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 >index 9058407c9ae36ea7cc13d9a5eba86848e3579c32..0000000000000000000000000000000000000000 >--- a/LayoutTests/webgpu/whlsl-compute.html >+++ /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>
<b>You cannot view the attachment while viewing its details because your browser does not support IFRAMEs. <a href="attachment.cgi?id=372073">View the attachment on a separate page</a>.</b>
View Attachment As Diff
View Attachment As Raw
Actions:
View
|
Formatted Diff
|
Diff
Attachments on
bug 198837
: 372073