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);