[WHLSL] Hook up compute
https://bugs.webkit.org/show_bug.cgi?id=198644
Reviewed by Saam Barati.
Source/WebCore:
This patch hooks up compute shaders in exactly the same way that vertex and fragment shaders
are hooked up. I've modified the two patchs (compute and rendering) to be almost exactly the
same code.
This patch also adds support for the WHLSL compiler to determine what the numthreads()
attribute in the shader says so that it can be hooked up to Metal's threads-per-threadgroup
argument in the dispatch call. There is some logic to make sure that there aren't two
numthreads() attributes on the same compute shader.
It also adds a little bit of type renaming. For built-in variables, sometimes Metal's type
doesn't always match WHLSL's (and HLSL's type). For example, in WHLSL and HLSL, SV_DispatchThreadID variables have to be a float3, but in Metal, they are a uint3.
Therefore, I've added a little bit of code during each entry point's pack and unpack stages
to handle this type conversion.
Test: webgpu/whlsl-compute.html
* Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp:
(WebCore::WHLSL::Metal::internalTypeForSemantic): Determine which Metal type corresponds to
each built-in variable.
(WebCore::WHLSL::Metal::EntryPointScaffolding::builtInsSignature): Perform the type
conversion.
(WebCore::WHLSL::Metal::EntryPointScaffolding::unpackResourcesAndNamedBuiltIns): Ditto.
(WebCore::WHLSL::Metal::VertexEntryPointScaffolding::VertexEntryPointScaffolding): Ditto.
(WebCore::WHLSL::Metal::VertexEntryPointScaffolding::helperTypes): Ditto.
(WebCore::WHLSL::Metal::VertexEntryPointScaffolding::pack): Ditto.
(WebCore::WHLSL::Metal::FragmentEntryPointScaffolding::FragmentEntryPointScaffolding): Ditto.
(WebCore::WHLSL::Metal::FragmentEntryPointScaffolding::helperTypes): Ditto.
(WebCore::WHLSL::Metal::FragmentEntryPointScaffolding::pack): Ditto.
(WebCore::WHLSL::Metal::ComputeEntryPointScaffolding::signature): Ditto.
* Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.h:
* Modules/webgpu/WHLSL/WHLSLComputeDimensions.cpp: Added. Add a pass to determine whether
or not any entry point has duplicate numthreads() attribute, and to determine what the
appropriate numthreads() values should be for the current entry point.
(WebCore::WHLSL::ComputeDimensionsVisitor::ComputeDimensionsVisitor):
(WebCore::WHLSL::ComputeDimensionsVisitor::computeDimensions const):
(WebCore::WHLSL::computeDimensions):
* Modules/webgpu/WHLSL/WHLSLComputeDimensions.h: Copied from Source/WebCore/Modules/webgpu/WHLSL/WHLSLPrepare.h.
* Modules/webgpu/WHLSL/WHLSLGatherEntryPointItems.cpp:
(WebCore::WHLSL::gatherEntryPointItems): Compute shaders don't need to have a semantic for their return type.
* Modules/webgpu/WHLSL/WHLSLPrepare.cpp:
(WebCore::WHLSL::prepare): Run the computeDimensions() pass.
* Modules/webgpu/WHLSL/WHLSLPrepare.h:
* Modules/webgpu/WHLSL/WHLSLPropertyResolver.cpp: In a left-value propertyAccessExpression,
the index expression can be a right-value. Treat it as such.
(WebCore::WHLSL::LeftValueSimplifier::finishVisiting):
(WebCore::WHLSL::LeftValueSimplifier::visit):
* Modules/webgpu/WHLSL/WHLSLStandardLibrary.txt: We need support for multiplication (for a
test) and float3 for SV_DispatchThreadID.
* Sources.txt:
* SourcesCocoa.txt:
* WebCore.xcodeproj/project.pbxproj:
* platform/graphics/gpu/GPUComputePipeline.h: Associate a compute dimensions with a particular
compute pipeline. This is how Metal knows what values to use for a dispatch.
(WebCore::GPUComputePipeline::computeDimensions const):
* platform/graphics/gpu/cocoa/GPUComputePassEncoderMetal.mm: Use the saved compute dimensions.
(WebCore::GPUComputePassEncoder::dispatch):
* platform/graphics/gpu/cocoa/GPUComputePipelineMetal.mm: Make the code match GPURenderPipelineMetal.
(WebCore::trySetMetalFunctions):
(WebCore::trySetFunctions):
(WebCore::convertComputePipelineDescriptor):
(WebCore::tryCreateMTLComputePipelineState):
(WebCore::GPUComputePipeline::tryCreate):
(WebCore::GPUComputePipeline::GPUComputePipeline):
(WebCore::tryCreateMtlComputeFunction): Deleted.
* platform/graphics/gpu/cocoa/GPUPipelineMetalConvertLayout.cpp: Added. Moved shared helper
functions to a file where they can be accessed by multiple places.
(WebCore::convertShaderStageFlags):
(WebCore::convertBindingType):
(WebCore::convertLayout):
* platform/graphics/gpu/cocoa/GPUPipelineMetalConvertLayout.h: Copied from Source/WebCore/Modules/webgpu/WHLSL/WHLSLPrepare.h.
* platform/graphics/gpu/cocoa/GPURenderPipelineMetal.mm: Delete the functions that were moved to GPUPipelineMetalConvertLayout.
(WebCore::trySetFunctions):
(WebCore::tryCreateMtlRenderPipelineState):
(WebCore::convertShaderStageFlags): Deleted.
(WebCore::convertBindingType): Deleted.
(WebCore::convertLayout): Deleted.
LayoutTests:
This doesn't thoroughly test compute, but it's at least enough to unblock the WHLSL testing effort.
* webgpu/compute-squares-expected.txt: Deleted. Covered by webgpu/whlsl-compute.html.
* webgpu/compute-squares.html: Deleted. Ditto.
* webgpu/whlsl-compute-expected.txt: Added.
* webgpu/whlsl-compute.html: Added.
git-svn-id: http://svn.webkit.org/repository/webkit/trunk@246396 268f45cc-cd09-0410-ab3c-d52691b4dbfc
diff --git a/LayoutTests/ChangeLog b/LayoutTests/ChangeLog
index dd60667..7abeb70 100644
--- a/LayoutTests/ChangeLog
+++ b/LayoutTests/ChangeLog
@@ -1,5 +1,19 @@
2019-06-12 Myles C. Maxfield <mmaxfield@apple.com>
+ [WHLSL] Hook up compute
+ https://bugs.webkit.org/show_bug.cgi?id=198644
+
+ Reviewed by Saam Barati.
+
+ This doesn't thoroughly test compute, but it's at least enough to unblock the WHLSL testing effort.
+
+ * webgpu/compute-squares-expected.txt: Deleted. Covered by webgpu/whlsl-compute.html.
+ * webgpu/compute-squares.html: Deleted. Ditto.
+ * webgpu/whlsl-compute-expected.txt: Added.
+ * webgpu/whlsl-compute.html: Added.
+
+2019-06-12 Myles C. Maxfield <mmaxfield@apple.com>
+
[WHLSL] Implement array references
https://bugs.webkit.org/show_bug.cgi?id=198163
diff --git a/LayoutTests/webgpu/compute-squares-expected.txt b/LayoutTests/webgpu/compute-squares-expected.txt
deleted file mode 100644
index 8b13789..0000000
--- a/LayoutTests/webgpu/compute-squares-expected.txt
+++ /dev/null
@@ -1 +0,0 @@
-
diff --git a/LayoutTests/webgpu/compute-squares.html b/LayoutTests/webgpu/compute-squares.html
deleted file mode 100644
index 39bffc0..0000000
--- a/LayoutTests/webgpu/compute-squares.html
+++ /dev/null
@@ -1,78 +0,0 @@
-<!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
new file mode 100644
index 0000000..6e1668a
--- /dev/null
+++ b/LayoutTests/webgpu/whlsl-compute-expected.txt
@@ -0,0 +1,12 @@
+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
new file mode 100644
index 0000000..9058407
--- /dev/null
+++ b/LayoutTests/webgpu/whlsl-compute.html
@@ -0,0 +1,88 @@
+<!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>
diff --git a/Source/WebCore/ChangeLog b/Source/WebCore/ChangeLog
index 06ae4ed..234d6f6 100644
--- a/Source/WebCore/ChangeLog
+++ b/Source/WebCore/ChangeLog
@@ -1,5 +1,89 @@
2019-06-12 Myles C. Maxfield <mmaxfield@apple.com>
+ [WHLSL] Hook up compute
+ https://bugs.webkit.org/show_bug.cgi?id=198644
+
+ Reviewed by Saam Barati.
+
+ This patch hooks up compute shaders in exactly the same way that vertex and fragment shaders
+ are hooked up. I've modified the two patchs (compute and rendering) to be almost exactly the
+ same code.
+
+ This patch also adds support for the WHLSL compiler to determine what the numthreads()
+ attribute in the shader says so that it can be hooked up to Metal's threads-per-threadgroup
+ argument in the dispatch call. There is some logic to make sure that there aren't two
+ numthreads() attributes on the same compute shader.
+
+ It also adds a little bit of type renaming. For built-in variables, sometimes Metal's type
+ doesn't always match WHLSL's (and HLSL's type). For example, in WHLSL and HLSL, SV_DispatchThreadID variables have to be a float3, but in Metal, they are a uint3.
+ Therefore, I've added a little bit of code during each entry point's pack and unpack stages
+ to handle this type conversion.
+
+ Test: webgpu/whlsl-compute.html
+
+ * Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp:
+ (WebCore::WHLSL::Metal::internalTypeForSemantic): Determine which Metal type corresponds to
+ each built-in variable.
+ (WebCore::WHLSL::Metal::EntryPointScaffolding::builtInsSignature): Perform the type
+ conversion.
+ (WebCore::WHLSL::Metal::EntryPointScaffolding::unpackResourcesAndNamedBuiltIns): Ditto.
+ (WebCore::WHLSL::Metal::VertexEntryPointScaffolding::VertexEntryPointScaffolding): Ditto.
+ (WebCore::WHLSL::Metal::VertexEntryPointScaffolding::helperTypes): Ditto.
+ (WebCore::WHLSL::Metal::VertexEntryPointScaffolding::pack): Ditto.
+ (WebCore::WHLSL::Metal::FragmentEntryPointScaffolding::FragmentEntryPointScaffolding): Ditto.
+ (WebCore::WHLSL::Metal::FragmentEntryPointScaffolding::helperTypes): Ditto.
+ (WebCore::WHLSL::Metal::FragmentEntryPointScaffolding::pack): Ditto.
+ (WebCore::WHLSL::Metal::ComputeEntryPointScaffolding::signature): Ditto.
+ * Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.h:
+ * Modules/webgpu/WHLSL/WHLSLComputeDimensions.cpp: Added. Add a pass to determine whether
+ or not any entry point has duplicate numthreads() attribute, and to determine what the
+ appropriate numthreads() values should be for the current entry point.
+ (WebCore::WHLSL::ComputeDimensionsVisitor::ComputeDimensionsVisitor):
+ (WebCore::WHLSL::ComputeDimensionsVisitor::computeDimensions const):
+ (WebCore::WHLSL::computeDimensions):
+ * Modules/webgpu/WHLSL/WHLSLComputeDimensions.h: Copied from Source/WebCore/Modules/webgpu/WHLSL/WHLSLPrepare.h.
+ * Modules/webgpu/WHLSL/WHLSLGatherEntryPointItems.cpp:
+ (WebCore::WHLSL::gatherEntryPointItems): Compute shaders don't need to have a semantic for their return type.
+ * Modules/webgpu/WHLSL/WHLSLPrepare.cpp:
+ (WebCore::WHLSL::prepare): Run the computeDimensions() pass.
+ * Modules/webgpu/WHLSL/WHLSLPrepare.h:
+ * Modules/webgpu/WHLSL/WHLSLPropertyResolver.cpp: In a left-value propertyAccessExpression,
+ the index expression can be a right-value. Treat it as such.
+ (WebCore::WHLSL::LeftValueSimplifier::finishVisiting):
+ (WebCore::WHLSL::LeftValueSimplifier::visit):
+ * Modules/webgpu/WHLSL/WHLSLStandardLibrary.txt: We need support for multiplication (for a
+ test) and float3 for SV_DispatchThreadID.
+ * Sources.txt:
+ * SourcesCocoa.txt:
+ * WebCore.xcodeproj/project.pbxproj:
+ * platform/graphics/gpu/GPUComputePipeline.h: Associate a compute dimensions with a particular
+ compute pipeline. This is how Metal knows what values to use for a dispatch.
+ (WebCore::GPUComputePipeline::computeDimensions const):
+ * platform/graphics/gpu/cocoa/GPUComputePassEncoderMetal.mm: Use the saved compute dimensions.
+ (WebCore::GPUComputePassEncoder::dispatch):
+ * platform/graphics/gpu/cocoa/GPUComputePipelineMetal.mm: Make the code match GPURenderPipelineMetal.
+ (WebCore::trySetMetalFunctions):
+ (WebCore::trySetFunctions):
+ (WebCore::convertComputePipelineDescriptor):
+ (WebCore::tryCreateMTLComputePipelineState):
+ (WebCore::GPUComputePipeline::tryCreate):
+ (WebCore::GPUComputePipeline::GPUComputePipeline):
+ (WebCore::tryCreateMtlComputeFunction): Deleted.
+ * platform/graphics/gpu/cocoa/GPUPipelineMetalConvertLayout.cpp: Added. Moved shared helper
+ functions to a file where they can be accessed by multiple places.
+ (WebCore::convertShaderStageFlags):
+ (WebCore::convertBindingType):
+ (WebCore::convertLayout):
+ * platform/graphics/gpu/cocoa/GPUPipelineMetalConvertLayout.h: Copied from Source/WebCore/Modules/webgpu/WHLSL/WHLSLPrepare.h.
+ * platform/graphics/gpu/cocoa/GPURenderPipelineMetal.mm: Delete the functions that were moved to GPUPipelineMetalConvertLayout.
+ (WebCore::trySetFunctions):
+ (WebCore::tryCreateMtlRenderPipelineState):
+ (WebCore::convertShaderStageFlags): Deleted.
+ (WebCore::convertBindingType): Deleted.
+ (WebCore::convertLayout): Deleted.
+
+2019-06-12 Myles C. Maxfield <mmaxfield@apple.com>
+
[WHLSL] Implement array references
https://bugs.webkit.org/show_bug.cgi?id=198163
diff --git a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp
index da8c057..3baa37f 100644
--- a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp
+++ b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp
@@ -178,6 +178,41 @@
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())
@@ -190,9 +225,11 @@
auto& namedBuiltIn = m_namedBuiltIns[i];
auto& item = m_entryPointItems.inputs[namedBuiltIn.indexInEntryPointItems];
auto& builtInSemantic = WTF::get<AST::BuiltInSemantic>(*item.semantic);
- auto mangledTypeName = m_typeNamer.mangledNameForType(*item.unnamedType);
+ auto internalType = internalTypeForSemantic(builtInSemantic);
+ if (internalType.isNull())
+ internalType = m_typeNamer.mangledNameForType(*item.unnamedType);
auto variableName = namedBuiltIn.variableName;
- stringBuilder.append(makeString(mangledTypeName, ' ', variableName, ' ', attributeForSemantic(builtInSemantic)));
+ stringBuilder.append(makeString(internalType, ' ', variableName, ' ', attributeForSemantic(builtInSemantic)));
}
return stringBuilder.toString();
}
@@ -299,9 +336,11 @@
}
for (auto& namedBuiltIn : m_namedBuiltIns) {
- auto& path = m_entryPointItems.inputs[namedBuiltIn.indexInEntryPointItems].path;
+ auto& item = m_entryPointItems.inputs[namedBuiltIn.indexInEntryPointItems];
+ auto& path = item.path;
auto& variableName = namedBuiltIn.variableName;
- stringBuilder.append(makeString(mangledInputPath(path), " = ", variableName, ";\n"));
+ auto mangledTypeName = m_typeNamer.mangledNameForType(*item.unnamedType);
+ stringBuilder.append(makeString(mangledInputPath(path), " = ", mangledTypeName, '(', variableName, ");\n"));
}
return stringBuilder.toString();
}
@@ -324,8 +363,13 @@
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));
}
}
@@ -346,10 +390,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 mangledTypeName = m_typeNamer.mangledNameForType(*outputItem.unnamedType);
+ auto& internalTypeName = m_namedOutputs[i].internalTypeName;
auto elementName = m_namedOutputs[i].elementName;
auto attribute = attributeForSemantic(*outputItem.semantic);
- stringBuilder.append(makeString(" ", mangledTypeName, ' ', elementName, ' ', attribute, ";\n"));
+ stringBuilder.append(makeString(" ", internalTypeName, ' ', elementName, ' ', attribute, ";\n"));
}
stringBuilder.append("};\n\n");
@@ -398,8 +442,9 @@
}
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, " = ", inputVariableName, mangledOutputPath(path), ";\n"));
+ stringBuilder.append(makeString(outputVariableName, '.', elementName, " = ", internalTypeName, '(', inputVariableName, mangledOutputPath(path), ");\n"));
}
return stringBuilder.toString();
}
@@ -424,8 +469,13 @@
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));
}
}
@@ -446,10 +496,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 mangledTypeName = m_typeNamer.mangledNameForType(*outputItem.unnamedType);
+ auto& internalTypeName = m_namedOutputs[i].internalTypeName;
auto elementName = m_namedOutputs[i].elementName;
auto attribute = attributeForSemantic(*outputItem.semantic);
- stringBuilder.append(makeString(" ", mangledTypeName, ' ', elementName, ' ', attribute, ";\n"));
+ stringBuilder.append(makeString(" ", internalTypeName, ' ', elementName, ' ', attribute, ";\n"));
}
stringBuilder.append("};\n\n");
@@ -498,8 +548,9 @@
}
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, " = ", inputVariableName, mangledOutputPath(path), ";\n"));
+ stringBuilder.append(makeString(outputVariableName, '.', elementName, " = ", internalTypeName, '(', inputVariableName, mangledOutputPath(path), ");\n"));
}
return stringBuilder.toString();
}
@@ -518,7 +569,7 @@
{
StringBuilder stringBuilder;
- stringBuilder.append(makeString("compute void ", functionName, '('));
+ stringBuilder.append(makeString("kernel 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 664f7fb..0f5cea5 100644
--- a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.h
+++ b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.h
@@ -129,6 +129,7 @@
struct NamedOutput {
String elementName;
+ String internalTypeName;
};
Vector<NamedOutput> m_namedOutputs;
};
@@ -157,6 +158,7 @@
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
new file mode 100644
index 0000000..c4b7d9c
--- /dev/null
+++ b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLComputeDimensions.cpp
@@ -0,0 +1,90 @@
+/*
+ * 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
new file mode 100644
index 0000000..de45081
--- /dev/null
+++ b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLComputeDimensions.h
@@ -0,0 +1,44 @@
+/*
+ * 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 996dd84..aa249c0 100644
--- a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLGatherEntryPointItems.cpp
+++ b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLGatherEntryPointItems.cpp
@@ -175,7 +175,8 @@
return WTF::nullopt;
}
Gatherer outputGatherer(intrinsics, functionDefinition.semantic() ? &*functionDefinition.semantic() : nullptr);
- outputGatherer.checkErrorAndVisit(functionDefinition.type());
+ if (*functionDefinition.entryPointType() != AST::EntryPointType::Compute)
+ 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 42848c7..d6b2e74 100644
--- a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLPrepare.cpp
+++ b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLPrepare.cpp
@@ -32,6 +32,7 @@
#include "WHLSLAutoInitializeVariables.h"
#include "WHLSLCheckDuplicateFunctions.h"
#include "WHLSLChecker.h"
+#include "WHLSLComputeDimensions.h"
#include "WHLSLFunctionStageChecker.h"
#include "WHLSLHighZombieFinder.h"
#include "WHLSLLiteralTypeChecker.h"
@@ -170,12 +171,16 @@
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 5bd6df0..eb2021b 100644
--- a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLPrepare.h
+++ b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLPrepare.h
@@ -42,9 +42,16 @@
};
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 a93e037..f80da99 100644
--- a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLPropertyResolver.cpp
+++ b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLPropertyResolver.cpp
@@ -774,6 +774,8 @@
ASSERT(propertyAccessExpression.base().typeAnnotation().leftAddressSpace());
ASSERT(propertyAccessExpression.anderFunction());
+ Visitor::visit(propertyAccessExpression.base());
+
Lexer::Token origin = propertyAccessExpression.origin();
auto* anderFunction = propertyAccessExpression.anderFunction();
@@ -804,7 +806,6 @@
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 1c4527a..5170e1a 100644
--- a/Source/WebCore/Modules/webgpu/WHLSL/WHLSLStandardLibrary.txt
+++ b/Source/WebCore/Modules/webgpu/WHLSL/WHLSLStandardLibrary.txt
@@ -432,15 +432,6 @@
native operator float(int);
native operator float(half);
-native float operator.x(float4);
-native float operator.y(float4);
-native float operator.z(float4);
-native float operator.w(float4);
-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);
@@ -448,6 +439,170 @@
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 ddx(float);
native float ddy(float);
diff --git a/Source/WebCore/Sources.txt b/Source/WebCore/Sources.txt
index 0d381d3..1b2cfd1 100644
--- a/Source/WebCore/Sources.txt
+++ b/Source/WebCore/Sources.txt
@@ -306,6 +306,7 @@
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 6048303..18bfd30 100644
--- a/Source/WebCore/SourcesCocoa.txt
+++ b/Source/WebCore/SourcesCocoa.txt
@@ -326,6 +326,7 @@
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 43007e6..aacba4c 100644
--- a/Source/WebCore/WebCore.xcodeproj/project.pbxproj
+++ b/Source/WebCore/WebCore.xcodeproj/project.pbxproj
@@ -6394,6 +6394,10 @@
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>"; };
@@ -25448,6 +25452,8 @@
C234A9AE21E92C1A003C984D /* WHLSLCheckDuplicateFunctions.h */,
1C840B9B21EC400900D0500D /* WHLSLChecker.cpp */,
1C840B9721EC400700D0500D /* WHLSLChecker.h */,
+ 1C86CA4B22AA19FF001BF961 /* WHLSLComputeDimensions.cpp */,
+ 1C86CA4C22AA19FF001BF961 /* WHLSLComputeDimensions.h */,
1CA0C2E421EED12A00A11860 /* WHLSLFunctionStageChecker.cpp */,
1CA0C2E521EED12A00A11860 /* WHLSLFunctionStageChecker.h */,
1C840B9921EC400800D0500D /* WHLSLGatherEntryPointItems.cpp */,
@@ -25987,6 +25993,8 @@
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 0359f29..65367c9 100644
--- a/Source/WebCore/platform/graphics/gpu/GPUComputePipeline.h
+++ b/Source/WebCore/platform/graphics/gpu/GPUComputePipeline.h
@@ -27,6 +27,7 @@
#if ENABLE(WEBGPU)
+#include "WHLSLPrepare.h"
#include <wtf/RefCounted.h>
#include <wtf/RefPtr.h>
#include <wtf/RetainPtr.h>
@@ -48,10 +49,13 @@
const PlatformComputePipeline* platformComputePipeline() const { return m_platformComputePipeline.get(); }
+ WHLSL::ComputeDimensions computeDimensions() const { return m_computeDimensions; }
+
private:
- GPUComputePipeline(PlatformComputePipelineSmartPtr&&);
+ GPUComputePipeline(PlatformComputePipelineSmartPtr&&, WHLSL::ComputeDimensions);
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 713c1b7..6bd28f7 100644
--- a/Source/WebCore/platform/graphics/gpu/cocoa/GPUComputePassEncoderMetal.mm
+++ b/Source/WebCore/platform/graphics/gpu/cocoa/GPUComputePassEncoderMetal.mm
@@ -97,11 +97,7 @@
BEGIN_BLOCK_OBJC_EXCEPTIONS;
- 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 threadsPerThreadgroup = MTLSizeMake(m_pipeline->computeDimensions().width, m_pipeline->computeDimensions().height, m_pipeline->computeDimensions().depth);
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 fbe43f4..d6eced4 100644
--- a/Source/WebCore/platform/graphics/gpu/cocoa/GPUComputePipelineMetal.mm
+++ b/Source/WebCore/platform/graphics/gpu/cocoa/GPUComputePipelineMetal.mm
@@ -30,68 +30,172 @@
#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 RetainPtr<MTLFunction> tryCreateMtlComputeFunction(const GPUPipelineStageDescriptor& stage)
+static bool trySetMetalFunctions(const char* const functionName, MTLLibrary *computeMetalLibrary, MTLComputePipelineDescriptor *mtlDescriptor, const String& computeEntryPointName)
{
- if (!stage.module->platformShaderModule() || stage.entryPoint.isNull()) {
- LOG(WebGPU, "GPUComputePipeline::tryCreate(): Invalid GPUShaderModule!");
- return nullptr;
- }
-
- RetainPtr<MTLFunction> function;
+#if LOG_DISABLED
+ UNUSED_PARAM(functionName);
+#endif
BEGIN_BLOCK_OBJC_EXCEPTIONS;
- function = adoptNS([stage.module->platformShaderModule() newFunctionWithName:stage.entryPoint]);
+
+ 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;
- if (!function)
- LOG(WebGPU, "GPUComputePipeline::tryCreate(): Cannot create compute MTLFunction \"%s\"!", stage.entryPoint.utf8().data());
-
- return function;
+ return false;
}
-static RetainPtr<MTLComputePipelineState> tryCreateMTLComputePipelineState(const GPUDevice& device, const GPUComputePipelineDescriptor& descriptor)
+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]);
+
+ 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 }};
+
+ return WTF::nullopt;
+}
+
+struct CreateResult {
+ RetainPtr<MTLComputePipelineState> pipelineState;
+ WHLSL::ComputeDimensions computeDimensions;
+};
+static Optional<CreateResult> tryCreateMTLComputePipelineState(const char* const functionName, const GPUDevice& device, const GPUComputePipelineDescriptor& descriptor)
{
if (!device.platformDevice()) {
LOG(WebGPU, "GPUComputePipeline::tryCreate(): Invalid GPUDevice!");
- return nullptr;
+ return WTF::nullopt;
}
- auto computeFunction = tryCreateMtlComputeFunction(descriptor.computeStage);
- if (!computeFunction)
- return nullptr;
+ auto convertResult = convertComputePipelineDescriptor(functionName, descriptor, device);
+ if (!convertResult)
+ return WTF::nullopt;
+ ASSERT(convertResult->pipelineDescriptor);
+ auto mtlDescriptor = convertResult->pipelineDescriptor;
- RetainPtr<MTLComputePipelineState> pipelineState;
- NSError *error = nil;
+ RetainPtr<MTLComputePipelineState> pipeline;
BEGIN_BLOCK_OBJC_EXCEPTIONS;
- pipelineState = adoptNS([device.platformDevice() newComputePipelineStateWithFunction:computeFunction.get() error:&error]);
+
+ NSError *error = nil;
+ pipeline = adoptNS([device.platformDevice() newComputePipelineStateWithDescriptor:mtlDescriptor.get() options:MTLPipelineOptionNone reflection:nil error:&error]);
+ if (!pipeline) {
+ LOG(WebGPU, "GPUComputePipeline::tryCreate(): %s!", error ? error.localizedDescription.UTF8String : "Unable to create MTLComputePipelineState!");
+ return WTF::nullopt;
+ }
+
END_BLOCK_OBJC_EXCEPTIONS;
- if (!pipelineState)
- LOG(WebGPU, "GPUComputePipeline::tryCreate(): %s!", error ? error.localizedDescription.UTF8String : "Unable to create MTLComputePipelineState!");
-
- return pipelineState;
+ return {{ pipeline, convertResult->computeDimensions }};
}
RefPtr<GPUComputePipeline> GPUComputePipeline::tryCreate(const GPUDevice& device, const GPUComputePipelineDescriptor& descriptor)
{
- auto mtlPipeline = tryCreateMTLComputePipelineState(device, descriptor);
- if (!mtlPipeline)
+ const char* const functionName = "GPURenderPipeline::create()";
+
+ auto createResult = tryCreateMTLComputePipelineState(functionName, device, descriptor);
+ if (!createResult)
return nullptr;
- return adoptRef(new GPUComputePipeline(WTFMove(mtlPipeline)));
+ return adoptRef(new GPUComputePipeline(WTFMove(createResult->pipelineState), createResult->computeDimensions));
}
-GPUComputePipeline::GPUComputePipeline(RetainPtr<MTLComputePipelineState>&& pipeline)
+GPUComputePipeline::GPUComputePipeline(RetainPtr<MTLComputePipelineState>&& pipeline, WHLSL::ComputeDimensions computeDimensions)
: 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
new file mode 100644
index 0000000..bd57688
--- /dev/null
+++ b/Source/WebCore/platform/graphics/gpu/cocoa/GPUPipelineMetalConvertLayout.cpp
@@ -0,0 +1,94 @@
+/*
+ * 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
new file mode 100644
index 0000000..8e29d93
--- /dev/null
+++ b/Source/WebCore/platform/graphics/gpu/cocoa/GPUPipelineMetalConvertLayout.h
@@ -0,0 +1,41 @@
+/*
+ * 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 4a9eebb..1745df0 100644
--- a/Source/WebCore/platform/graphics/gpu/cocoa/GPURenderPipelineMetal.mm
+++ b/Source/WebCore/platform/graphics/gpu/cocoa/GPURenderPipelineMetal.mm
@@ -30,6 +30,7 @@
#import "GPUDevice.h"
#import "GPULimits.h"
+#import "GPUPipelineMetalConvertLayout.h"
#import "GPUUtils.h"
#import "Logging.h"
#import "WHLSLPrepare.h"
@@ -98,35 +99,6 @@
}
}
-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) {
@@ -368,34 +340,6 @@
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
@@ -466,8 +410,6 @@
if (!whlslCompileResult)
return false;
- WTFLogAlways("Metal Source: %s", whlslCompileResult->metalSource.utf8().data());
-
NSError *error = nil;
BEGIN_BLOCK_OBJC_EXCEPTIONS;
@@ -542,6 +484,11 @@
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;
@@ -550,7 +497,7 @@
BEGIN_BLOCK_OBJC_EXCEPTIONS;
- NSError *error = [NSError errorWithDomain:@"com.apple.WebKit.GPU" code:1 userInfo:nil];
+ NSError *error = nil;
pipeline = adoptNS([device.platformDevice() newRenderPipelineStateWithDescriptor:mtlDescriptor.get() error:&error]);
if (!pipeline)
LOG(WebGPU, "%s: %s!", functionName, error.localizedDescription.UTF8String);
