Diff
Modified: trunk/LayoutTests/ChangeLog (246627 => 246628)
--- trunk/LayoutTests/ChangeLog 2019-06-20 02:18:21 UTC (rev 246627)
+++ trunk/LayoutTests/ChangeLog 2019-06-20 03:18:37 UTC (rev 246628)
@@ -1,3 +1,47 @@
+2019-06-19 Justin Fan <justin_...@apple.com>
+
+ [WHLSL] Create a shading language test harness
+ https://bugs.webkit.org/show_bug.cgi?id=198978
+
+ Reviewed by Myles C. Maxfield.
+
+ Introduce a test harness that can be used to test WebGPU shader compilation and functionality.
+ Currently using MSL.
+ Will be replaced with WHLSL as it gains the minimum features needed to support.
+
+ * webgpu/js/whlsl-test-harness.js: Added.
+ (isVectorType):
+ (convertTypeToArrayType):
+ (convertTypeToWHLSLType):
+ (Data):
+ (Data.prototype.async.getArrayBuffer):
+ (Data.prototype.get type):
+ (Data.prototype.get isPointer):
+ (Data.prototype.get buffer):
+ (Data.prototype.get byteLength):
+ (Harness.prototype._initialize):
+ (Harness.prototype.async.callTypedFunction):
+ (Harness.prototype.async.callVoidFunction):
+ (Harness.prototype._setUpArguments):
+ (Harness.prototype._callFunction):
+ (Harness):
+ (harness._initialize.async):
+ (makeBool):
+ (makeInt):
+ (makeUchar):
+ (makeUint):
+ (makeFloat):
+ (makeFloat4):
+ (async.callBoolFunction):
+ (async.callIntFunction):
+ (async.callUcharFunction):
+ (async.callUintFunction):
+ (async.callFloatFunction):
+ (async.callFloat4Function):
+ (callVoidFunction):
+ * webgpu/whlsl-harness-test-expected.txt: Added.
+ * webgpu/whlsl-harness-test.html: Added.
+
2019-06-19 Saam Barati <sbar...@apple.com>
[WHLSL] The checker needs to resolve types for the anonymous variables in ReadModifyWrite expressions
Added: trunk/LayoutTests/webgpu/js/whlsl-test-harness.js (0 => 246628)
--- trunk/LayoutTests/webgpu/js/whlsl-test-harness.js (rev 0)
+++ trunk/LayoutTests/webgpu/js/whlsl-test-harness.js 2019-06-20 03:18:37 UTC (rev 246628)
@@ -0,0 +1,386 @@
+/* Type Utilties */
+
+// FIXME: Support all WHLSL scalar and vector types.
+// FIXME: Support textures and samplers.
+const Types = Object.freeze({
+ BOOL: Symbol("bool"),
+ INT: Symbol("int"),
+ UCHAR: Symbol("uchar"),
+ UINT: Symbol("uint"),
+ FLOAT: Symbol("float"),
+ FLOAT4: Symbol("float4"),
+ MAX_SIZE: 16 // This needs to be big enough to hold any singular WHLSL type.
+});
+
+function isVectorType(type)
+{
+ switch(type) {
+ case Types.FLOAT4:
+ return true;
+ default:
+ return false;
+ }
+}
+
+function convertTypeToArrayType(type)
+{
+ switch(type) {
+ case Types.BOOL:
+ return Uint8Array;
+ case Types.INT:
+ return Int32Array;
+ case Types.UCHAR:
+ return Uint8Array;
+ case Types.UINT:
+ return Uint32Array;
+ case Types.FLOAT:
+ case Types.FLOAT4:
+ return Float32Array;
+ default:
+ throw new Error("Invalid TYPE provided!");
+ }
+}
+
+function convertTypeToWHLSLType(type)
+{
+ switch(type) {
+ case Types.BOOL:
+ return "bool";
+ case Types.INT:
+ return "int";
+ case Types.UCHAR:
+ return "uchar";
+ case Types.UINT:
+ return "uint";
+ case Types.FLOAT:
+ return "float";
+ case Types.FLOAT4:
+ return "float4";
+ default:
+ throw new Error("Invalid TYPE provided!");
+ }
+}
+
+/* Harness Classes */
+
+class Data {
+ /**
+ * Upload typed data to and return a wrapper of a GPUBuffer.
+ * @param {Types} type - The WHLSL type to be stored in this Data.
+ * @param {Number or Array[Number]} values - The raw data to be uploaded.
+ */
+ constructor(harness, type, values, isPointer = false)
+ {
+ // One or more scalars in an array can be accessed through a pointer to buffer.
+ // However, vector types are also created via an array of scalars.
+ // This ensures that buffers of just one vector are usable in a test function.
+ if (Array.isArray(values))
+ this._isPointer = isVectorType(type) ? isPointer : true;
+ else {
+ this._isPointer = false;
+ values = [values];
+ }
+
+ this._type = type;
+ this._byteLength = (convertTypeToArrayType(type)).BYTES_PER_ELEMENT * values.length;
+
+ const [buffer, arrayBuffer] = harness._device.createBufferMapped({
+ size: this._byteLength,
+ usage: GPUBufferUsage.STORAGE | GPUBufferUsage.MAP_READ
+ });
+
+ const typedArray = new (convertTypeToArrayType(type))(arrayBuffer);
+ typedArray.set(values);
+ buffer.unmap();
+
+ this._buffer = buffer;
+ }
+
+ /**
+ * @returns An ArrayBuffer containing the contents of this Data.
+ */
+ async getArrayBuffer()
+ {
+ let result;
+ try {
+ result = await this._buffer.mapReadAsync();
+ this._buffer.unmap();
+ } catch {
+ throw new Error("Data error: Unable to get ArrayBuffer!");
+ }
+ return result;
+ }
+
+ get type() { return this._type; }
+ get isPointer() { return this._isPointer; }
+ get buffer() { return this._buffer; }
+ get byteLength() { return this._byteLength; }
+}
+
+class Harness {
+ constructor()
+ {
+ this._shaderHeader = `#include <metal_stdlib>
+ using namespace metal;
+ `;
+ }
+
+ _initialize(callback)
+ {
+ callback.bind(this)();
+ }
+
+ /**
+ * Return the return value of a WHLSL function.
+ * @param {Types} type - The return type of the WHLSL function.
+ * @param {String} functions - Custom WHLSL code to be tested.
+ * @param {String} name - The name of the WHLSL function which must be present in 'functions'.
+ * @param {Data or Array[Data]} args - Data arguments to be passed to the call of 'name'.
+ * @returns {TypedArray} - A typed array containing the return value of the function call.
+ */
+ async callTypedFunction(type, functions, name, args)
+ {
+ const [argsLayouts, argsResourceBindings, argsStructCode, functionCallArgs] = this._setUpArguments(args);
+
+ if (!this._resultBuffer) {
+ this._resultBuffer = this._device.createBuffer({
+ size: Types.MAX_SIZE,
+ usage: GPUBufferUsage.STORAGE | GPUBufferUsage.MAP_READ
+ });
+ }
+
+ argsLayouts.unshift({
+ binding: 0,
+ visibility: GPUShaderStageBit.COMPUTE,
+ type: "storage-buffer"
+ });
+ argsResourceBindings.unshift({
+ binding: 0,
+ resource: {
+ buffer: this._resultBuffer,
+ size: Types.MAX_SIZE
+ }
+ });
+
+ const code = this._shaderHeader + functions + `
+ struct _compute_args {
+ device ${convertTypeToWHLSLType(type)}* result [[id(0)]];
+ ${argsStructCode}};
+
+ kernel void _compute_main(device _compute_args& args [[buffer(0)]])
+ {
+ *args.result = ${name}(${functionCallArgs.join(", ")});
+ }
+ `;
+
+ this._callFunction(code, argsLayouts, argsResourceBindings);
+
+ try {
+ var result = await this._resultBuffer.mapReadAsync();
+ } catch {
+ throw new Error("Harness error: Unable to read results!");
+ }
+ const array = new (convertTypeToArrayType(type))(result);
+ this._resultBuffer.unmap();
+
+ return array;
+ }
+
+ /**
+ * Call a WHLSL function to modify the value of argument(buffer)s.
+ * @param {String} functions - Custom WHLSL code to be tested.
+ * @param {String} name - The name of the WHLSL function which must be present in 'functions'.
+ * @param {Data or Array[Data]} args - Data arguments to be passed to the call of 'name'.
+ */
+ async callVoidFunction(functions, name, args)
+ {
+ const [argsLayouts, argsResourceBindings, argsStructCode, functionCallArgs] = this._setUpArguments(args);
+
+ const code = this._shaderHeader + functions + `
+ struct _compute_args {
+ ${argsStructCode}};
+
+ kernel void _compute_main(device _compute_args& args [[buffer(0)]])
+ {
+ ${name}(${functionCallArgs.join(", ")});
+ }
+ `;
+
+ this._callFunction(code, argsLayouts, argsResourceBindings);
+ }
+
+ _setUpArguments(args)
+ {
+ if (!Array.isArray(args)) {
+ if (args instanceof Data)
+ args = [args];
+ else if (!args)
+ args = [];
+ }
+
+ // FIXME: Replace with WHLSL.
+ // Expand bind group structure to represent any arguments.
+ let argsStructCode = "";
+ let functionCallArgs = [];
+ let argsLayouts = [];
+ let argsResourceBindings = [];
+
+ for (let i = 1; i <= args.length; ++i) {
+ const arg = args[i - 1];
+ argsStructCode += `device ${convertTypeToWHLSLType(arg.type)}* arg${i} [[id(${i})]];
+ `;
+ const optionalDeref = (!arg.isPointer) ? "*" : "";
+ functionCallArgs.push(optionalDeref + `args.arg${i}`);
+ argsLayouts.push({
+ binding: i,
+ visibility: GPUShaderStageBit.COMPUTE,
+ type: "storage-buffer"
+ });
+ argsResourceBindings.push({
+ binding: i,
+ resource: {
+ buffer: arg.buffer,
+ size: arg.byteLength
+ }
+ });
+ }
+
+ return [argsLayouts, argsResourceBindings, argsStructCode, functionCallArgs];
+ }
+
+ _callFunction(code, argsLayouts, argsResourceBindings)
+ {
+ const shaders = this._device.createShaderModule({ code: code });
+ // FIXME: Compile errors should be caught and reported here.
+ const pipeline = this._device.createComputePipeline({
+ computeStage: {
+ module: shaders,
+ entryPoint: "_compute_main"
+ }
+ });
+
+ const layout = this._device.createBindGroupLayout({
+ bindings: argsLayouts
+ });
+
+ const bindGroup = this._device.createBindGroup({
+ layout: layout,
+ bindings: argsResourceBindings
+ });
+
+ const commandEncoder = this._device.createCommandEncoder();
+ const passEncoder = commandEncoder.beginComputePass();
+ passEncoder.setBindGroup(0, bindGroup);
+ passEncoder.setPipeline(pipeline);
+ passEncoder.dispatch(1, 1, 1);
+ passEncoder.endPass();
+
+ this._device.getQueue().submit([commandEncoder.finish()]);
+ }
+}
+
+/* Harness Setup */
+
+const harness = new Harness();
+harness._initialize(async () => {
+ try {
+ const adapter = await navigator.gpu.requestAdapter();
+ harness._device = await adapter.requestDevice();
+ } catch (e) {
+ throw new Error("Harness error: Unable to acquire GPUDevice!");
+ }
+});
+
+/* Global Helper Functions */
+
+/**
+ * The make___ functions are wrappers around the Data constructor.
+ * Values passed in as an array will be passed in via a device-addressed pointer type in the shader.
+ * @param {Boolean, Number, or Array} values - The data to be stored on the GPU.
+ * @returns A new Data object with storage allocated to store values.
+ */
+function makeBool(values)
+{
+ return new Data(harness, Types.BOOL, values);
+}
+
+function makeInt(values)
+{
+ return new Data(harness, Types.INT, values);
+}
+
+function makeUchar(values)
+{
+ return new Data(harness, Types.UCHAR, values);
+}
+
+function makeUint(values)
+{
+ return new Data(harness, Types.UINT, values);
+}
+
+function makeFloat(values)
+{
+ return new Data(harness, Types.FLOAT, values);
+}
+
+/**
+ * @param {Array or Array[Array]} values - 1D or 2D array of float values.
+ * The total number of float values must be divisible by 4.
+ * A single 4-element array of floats will be treated as a single float4 argument in the shader.
+ */
+function makeFloat4(values)
+{
+ const originalLength = values.length;
+ // This works because float4 is tightly packed.
+ // When implementing other vector types, add padding if needed.
+ values = values.flat();
+ if (values.length % 4 != 0)
+ throw new Error("makeFloat4: Invalid number of elements!");
+ return new Data(harness, Types.FLOAT4, values, originalLength === 1 || values.length > 4);
+}
+
+/**
+ * @param {String} functions - Shader source code that must contain a definition for 'name'.
+ * @param {String} name - The function to be called from 'functions'.
+ * @param {Data or Array[Data]} args - The arguments to be passed to the call of 'name'.
+ * @returns A Promise that resolves to the return value of a call to 'name' with 'args'.
+ */
+async function callBoolFunction(functions, name, args)
+{
+ return !!(await harness.callTypedFunction(Types.BOOL, functions, name, args))[0];
+}
+
+async function callIntFunction(functions, name, args)
+{
+ return (await harness.callTypedFunction(Types.INT, functions, name, args))[0];
+}
+
+async function callUcharFunction(functions, name, args)
+{
+ return (await harness.callTypedFunction(Types.UCHAR, functions, name, args))[0];
+}
+
+async function callUintFunction(functions, name, args)
+{
+ return (await harness.callTypedFunction(Types.UINT, functions, name, args))[0];
+}
+
+async function callFloatFunction(functions, name, args)
+{
+ return (await harness.callTypedFunction(Types.FLOAT, functions, name, args))[0];
+}
+
+async function callFloat4Function(functions, name, args)
+{
+ return (await harness.callTypedFunction(Types.FLOAT4, functions, name, args)).subarray(0, 4);
+}
+
+/**
+ * Does not return a Promise. To observe the results of a call,
+ * call 'getArrayBuffer' on the Data object retaining your output buffer.
+ */
+function callVoidFunction(functions, name, args)
+{
+ harness.callVoidFunction(functions, name, args);
+}
\ No newline at end of file
Added: trunk/LayoutTests/webgpu/whlsl-harness-test-expected.txt (0 => 246628)
--- trunk/LayoutTests/webgpu/whlsl-harness-test-expected.txt (rev 0)
+++ trunk/LayoutTests/webgpu/whlsl-harness-test-expected.txt 2019-06-20 03:18:37 UTC (rev 246628)
@@ -0,0 +1,35 @@
+
+PASS Return a literal of type bool.
+PASS Return an expected float4 value.
+PASS Return an expected int value.
+PASS Return an expected uchar value.
+PASS Return an expected uint value.
+PASS Return an expected float value.
+PASS Upload and return a bool value.
+PASS Return an expected float4 value.
+PASS Return an expected int value.
+PASS Return an expected uchar value.
+PASS Return an expected uint value.
+PASS Return an expected float value.
+PASS Upload many bool values and return a calculated result.
+PASS Return an expected float4 value.
+PASS Return an expected int value.
+PASS Return an expected uchar value.
+PASS Return an expected uint value.
+PASS Return an expected float value.
+PASS Access and return a single bool through a bool*.
+PASS Return an expected float4 value.
+PASS Return an expected int value.
+PASS Return an expected uchar value.
+PASS Return an expected uint value.
+PASS Return an expected float value.
+PASS Access multiple bools through various buffers and return a bool.
+PASS Return an expected float4 value.
+PASS Return an expected int value.
+PASS Return an expected uchar value.
+PASS Return an expected uint value.
+PASS Return an expected float value.
+PASS Upload and calculate a result from varied argument types.
+PASS Store into a float4*.
+PASS Upload a uchar* and store into a uchar*.
+
Added: trunk/LayoutTests/webgpu/whlsl-harness-test.html (0 => 246628)
--- trunk/LayoutTests/webgpu/whlsl-harness-test.html (rev 0)
+++ trunk/LayoutTests/webgpu/whlsl-harness-test.html 2019-06-20 03:18:37 UTC (rev 246628)
@@ -0,0 +1,208 @@
+<!DOCTYPE html><!-- webkit-test-runner [ experimental:WebGPUEnabled=true ] -->
+<html>
+<meta charset=utf-8>
+<title>Test the WHLSL test harness.</title>
+<script src=""
+<script src=""
+<script src=""
+<script>
+const epsilon = 0.0001;
+
+const numericScalarTypes = ["int", "uchar", "uint", "float"];
+
+const numericScalarFuncs = {
+ "int": callIntFunction,
+ "uchar": callUcharFunction,
+ "uint": callUintFunction,
+ "float": callFloatFunction
+};
+
+const scalarArgMakers = {
+ "int": makeInt,
+ "uchar": makeUchar,
+ "uint": makeUint,
+ "float": makeFloat
+};
+
+let tests = {};
+
+tests.literals = () => {
+ checkBools("Return a literal of type bool.", "return true;");
+ checkFloat4s("return float4(0, 1, 2, 3);");
+ checkNumericScalars("return 42;", [], 42);
+};
+
+tests.singleArgument = () => {
+ checkBools("Upload and return a bool value.", "return in0;", [true]);
+ checkFloat4s("return in0.wzyx;", [[3, 2, 1, 0]]);
+ checkNumericScalars("return in0;", [42], 42);
+};
+
+tests.manyArguments = () => {
+ checkBools("Upload many bool values and return a calculated result.",
+ "return in0 & in1 & in2 & in3 & in4 & in5 & in6 & in7;",
+ [true, true, true, true, true, true, true, true]);
+
+ const body = `return in0 + in1 + in2 + in3 + in4 + in5 + in6 + in7;`;
+ let args = [];
+ for (let i = 0; i < 8; ++i)
+ args.push([0, 1, 2, 3]);
+ checkFloat4s(body, args, [0, 8, 16, 24]);
+ checkNumericScalars(body, [0, 1, 2, 3, 4, 5, 6, 7], 28);
+};
+
+tests.buffersWithOneValue = () => {
+ const body = `return in0[0];`
+ checkBools("Access and return a single bool through a bool*.", body, [[true]]);
+ checkFloat4s(body, [[[0, 1, 2, 3]]]);
+ checkNumericScalars(body, [[42]], 42);
+};
+
+tests.multipleBufferArguments = () => {
+ checkBools("Access multiple bools through various buffers and return a bool.",
+ "return in0[0] & in0[1] & in0[2] & in1 & in2[0];",
+ [[true, true, true], true, [true]]);
+
+ const body = `return in0[0] + in0[1] + in0[2] + in1 + in2[0];`;
+ const vector = [0, 1, 2, 3];
+ checkFloat4s(body, [[vector, vector, vector], vector, [vector]], [0, 5, 10, 15]);
+ checkNumericScalars(body, [[0, 1, 2], 3, [4]], 10);
+};
+
+tests.multipleArgumentTypes = () => {
+ const src = "" test(int i, uchar c, device uint* u, bool b, device bool* bs, float4 f4, device float* fs)
+ {
+ if (b && bs[0] && bs[1])
+ return i + c + u[0] + f4.x + f4.y + f4.z + f4.w + fs[0] + fs[1];
+
+ return 0;
+ }`;
+ const i = makeInt(1);
+ const c = makeUchar(2);
+ const u = makeUint([3]);
+ const b = makeBool(true);
+ const bs = makeBool([true, true]);
+ const f4 = makeFloat4([4, 5, 6, 7]);
+ const fs = makeFloat([8, 9]);
+ promise_test(() => {
+ return callFloatFunction(src, "test", [i, c, u, b, bs, f4, fs]).then(result => {
+ assert_approx_equals(result, 45, epsilon, "Test returned expected value.");
+ });
+ }, "Upload and calculate a result from varied argument types.");
+};
+
+tests.bufferStores = () => {
+ let src = "" test(device float4* out) {
+ *out = float4(0, 1, 2, 3);
+ }`;
+ const float4Out = makeFloat4([[0, 0, 0, 0]]);
+ callVoidFunction(src, "test", float4Out);
+
+ promise_test(() => {
+ return float4Out.getArrayBuffer().then(arrayBuffer => {
+ const result = new Float32Array(arrayBuffer);
+ for (let i; i < 4; ++i) {
+ assert_approx_equals(result[i], i, "Test stored expected values.");
+ }
+ });
+ }, "Store into a float4*.");
+
+ src = "" test(device uchar* in, device uchar* out) {
+ for (uint i = 0; i < 5; ++i)
+ out[i] = in[i];
+ }`;
+ const array = [0, 1, 2, 3, 4];
+ const input = makeUchar(array);
+ const output = makeUchar([0, 0, 0, 0, 0]);
+ callVoidFunction(src, "test", [input, output]);
+
+ promise_test(() => {
+ return output.getArrayBuffer().then(arrayBuffer => {
+ const result = new Uint8Array(arrayBuffer);
+ assert_array_equals(array, result, "Test stored expected values.");
+ });
+ }, "Upload a uchar* and store into a uchar*.");
+};
+
+window.addEventListener("load", () => {
+ for (const name in tests) {
+ tests[name]();
+ }
+});
+
+/* Helper functions */
+
+const checkNumericScalars = (body, argValues, expected) => {
+ let functions = [];
+ let src = ""
+ for (let type of numericScalarTypes) {
+ const name = `${type}Test`;
+
+ let inputArgs = [];
+ let values = [];
+ for (let i = 0; i < argValues.length; ++i) {
+ const isPointer = Array.isArray(argValues[i]);
+ inputArgs.push(`${isPointer ? "device " : ""}${type}${isPointer ? "*" : ""} in${i}`);
+ values.push(scalarArgMakers[type](argValues[i]));
+ }
+
+ src += `${type} ${name}(${inputArgs.join(", ")}) { ${body} }
+ `;
+ functions.push({ type: type, name: name, args: values, expected: expected });
+ }
+
+ for (const f of functions) {
+ const callFunc = numericScalarFuncs[f.type];
+ promise_test(async () => {
+ return callFunc(src, f.name, f.args).then(result => {
+ assert_approx_equals(result, f.expected, epsilon, "Test returned expected value.");
+ });
+ }, `Return an expected ${f.type} value.`);
+ }
+};
+
+const checkBools = (msg = "Return an expected bool value.", body, argValues = [], expected = true) => {
+ let src = ""
+ let inputArgs = [];
+ let values = [];
+ for (let i = 0; i < argValues.length; ++i) {
+ const isPointer = Array.isArray(argValues[i]);
+ inputArgs.push(`${isPointer ? "device " : ""}bool${isPointer ? "*" : ""} in${i}`);
+ values.push(makeBool(argValues[i]));
+ }
+
+ src += `bool boolTest(${inputArgs.join(", ")}) { ${body} }
+ `;
+
+ promise_test(async () => {
+ return callBoolFunction(src, "boolTest", values).then(result => {
+ assert_equals(result, expected, "Test returned expected value.");
+ });
+ }, msg);
+};
+
+const checkFloat4s = (body, argValues = [], expected = [0, 1, 2, 3]) => {
+ let src = ""
+ let inputArgs = [];
+ let values = [];
+
+ for (let i = 0; i < argValues.length; ++i) {
+ // Support arrays of float4, including one with a single float4.
+ const totalLength = argValues[i].flat().length;
+ const isPointer = argValues[i].length === 1 || totalLength > 4;
+ inputArgs.push(`${isPointer ? "device " : ""}float4${isPointer ? "*" : ""} in${i}`);
+ values.push(makeFloat4(argValues[i]));
+ }
+
+ src += `float4 float4Test(${inputArgs.join(", ")}) { ${body} }
+ `;
+
+ promise_test(async () => {
+ return callFloat4Function(src, "float4Test", values).then(result => {
+ for (let i = 0; i < 4; ++i)
+ assert_approx_equals(result[i], expected[i], epsilon, "Test returned expected value.");
+ });
+ }, "Return an expected float4 value.");
+}
+</script>
+</html>
\ No newline at end of file
Modified: trunk/Source/WebCore/ChangeLog (246627 => 246628)
--- trunk/Source/WebCore/ChangeLog 2019-06-20 02:18:21 UTC (rev 246627)
+++ trunk/Source/WebCore/ChangeLog 2019-06-20 03:18:37 UTC (rev 246628)
@@ -1,3 +1,21 @@
+2019-06-19 Justin Fan <justin_...@apple.com>
+
+ [WHLSL] Create a shading language test harness
+ https://bugs.webkit.org/show_bug.cgi?id=198978
+
+ Reviewed by Myles C. Maxfield.
+
+ When creating MTLArgumentEncoders for argument buffers, the user's arguments
+ must match the order that they are declared in the shader. Move back-end information
+ such as buffer lengths to the end of the argument arrays.
+
+ Test: webgpu/whlsl-harness-test.html
+
+ * Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp:
+ (WebCore::WHLSL::Metal::EntryPointScaffolding::resourceHelperTypes):
+ * platform/graphics/gpu/cocoa/GPUBindGroupLayoutMetal.mm:
+ (WebCore::GPUBindGroupLayout::tryCreate):
+
2019-06-19 Saam Barati <sbar...@apple.com>
[WHLSL] The checker needs to resolve types for the anonymous variables in ReadModifyWrite expressions
Modified: trunk/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp (246627 => 246628)
--- trunk/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp 2019-06-20 02:18:21 UTC (rev 246627)
+++ trunk/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp 2019-06-20 03:18:37 UTC (rev 246628)
@@ -37,6 +37,7 @@
#include "WHLSLStageInOutSemantic.h"
#include "WHLSLStructureDefinition.h"
#include "WHLSLTypeNamer.h"
+#include <algorithm>
#include <wtf/Optional.h>
#include <wtf/text/StringBuilder.h>
#include <wtf/text/StringConcatenateNumbers.h>
@@ -143,6 +144,7 @@
StringBuilder stringBuilder;
for (size_t i = 0; i < m_layout.size(); ++i) {
stringBuilder.append(makeString("struct ", m_namedBindGroups[i].structName, " {\n"));
+ Vector<std::pair<unsigned, String>> structItems;
for (size_t j = 0; j < m_layout[i].bindings.size(); ++j) {
auto iterator = m_resourceMap.find(&m_layout[i].bindings[j]);
if (iterator == m_resourceMap.end())
@@ -153,10 +155,15 @@
auto addressSpace = toString(referenceType.addressSpace());
auto elementName = m_namedBindGroups[i].namedBindings[j].elementName;
auto index = m_namedBindGroups[i].namedBindings[j].index;
- stringBuilder.append(makeString(" ", addressSpace, " ", mangledTypeName, "* ", elementName, " [[id(", index, ")]];\n"));
+ structItems.append(std::make_pair(index, makeString(" ", addressSpace, " ", mangledTypeName, "* ", elementName, " [[id(", index, ")]];\n")));
if (auto lengthInformation = m_namedBindGroups[i].namedBindings[j].lengthInformation)
- stringBuilder.append(makeString(" uint2 ", lengthInformation->elementName, " [[id(", lengthInformation->index, ")]];\n"));
+ structItems.append(std::make_pair(lengthInformation->index, makeString("uint2 ", lengthInformation->elementName, " [[id(", lengthInformation->index, ")]];")));
}
+ std::sort(structItems.begin(), structItems.end(), [](const std::pair<unsigned, String>& left, const std::pair<unsigned, String>& right) {
+ return left.first < right.first;
+ });
+ for (const auto& structItem : structItems)
+ stringBuilder.append(makeString(" ", structItem.second, '\n'));
stringBuilder.append("};\n\n");
}
return stringBuilder.toString();
Modified: trunk/Source/WebCore/platform/graphics/gpu/cocoa/GPUBindGroupLayoutMetal.mm (246627 => 246628)
--- trunk/Source/WebCore/platform/graphics/gpu/cocoa/GPUBindGroupLayoutMetal.mm 2019-06-20 02:18:21 UTC (rev 246627)
+++ trunk/Source/WebCore/platform/graphics/gpu/cocoa/GPUBindGroupLayoutMetal.mm 2019-06-20 03:18:37 UTC (rev 246628)
@@ -96,7 +96,7 @@
return nullptr;
}
- ArgumentArray vertexArgsArray, fragmentArgsArray, computeArgsArray;
+ ArgumentArray vertexArgs, fragmentArgs, computeArgs, vertexLengths, fragmentLengths, computeLengths;
BindingsMapType bindingsMap;
unsigned internalName = 0;
@@ -137,8 +137,8 @@
return nullptr;
}
- auto addIndices = [&](ArgumentArray& array) -> bool {
- appendArgumentToArray(array, mtlArgument);
+ auto addIndices = [&](ArgumentArray& args, ArgumentArray& lengths) -> bool {
+ appendArgumentToArray(args, mtlArgument);
if (extraIndex) {
RetainPtr<MTLArgumentDescriptor> mtlArgument = argumentDescriptor(MTLDataTypeUInt2, *extraIndex);
if (!mtlArgument) {
@@ -145,25 +145,31 @@
LOG(WebGPU, "GPUBindGroupLayout::tryCreate(): Unable to create MTLArgumentDescriptor for binding %u!", binding.binding);
return false;
}
- appendArgumentToArray(array, mtlArgument);
+ appendArgumentToArray(lengths, mtlArgument);
}
return true;
};
- if ((binding.visibility & GPUShaderStageBit::Flags::Vertex) && !addIndices(vertexArgsArray))
+ if ((binding.visibility & GPUShaderStageBit::Flags::Vertex) && !addIndices(vertexArgs, vertexLengths))
return nullptr;
- if ((binding.visibility & GPUShaderStageBit::Flags::Fragment) && !addIndices(fragmentArgsArray))
+ if ((binding.visibility & GPUShaderStageBit::Flags::Fragment) && !addIndices(fragmentArgs, fragmentLengths))
return nullptr;
- if ((binding.visibility & GPUShaderStageBit::Flags::Compute) && !addIndices(computeArgsArray))
+ if ((binding.visibility & GPUShaderStageBit::Flags::Compute) && !addIndices(computeArgs, computeLengths))
return nullptr;
}
+ BEGIN_BLOCK_OBJC_EXCEPTIONS;
+ [vertexArgs addObjectsFromArray:vertexLengths.get()];
+ [fragmentArgs addObjectsFromArray:fragmentLengths.get()];
+ [computeArgs addObjectsFromArray:computeLengths.get()];
+ END_BLOCK_OBJC_EXCEPTIONS;
+
RetainPtr<MTLArgumentEncoder> vertex, fragment, compute;
- if (vertexArgsArray && !(vertex = tryCreateMtlArgumentEncoder(device, vertexArgsArray)))
+ if (vertexArgs && !(vertex = tryCreateMtlArgumentEncoder(device, vertexArgs)))
return nullptr;
- if (fragmentArgsArray && !(fragment = tryCreateMtlArgumentEncoder(device, fragmentArgsArray)))
+ if (fragmentArgs && !(fragment = tryCreateMtlArgumentEncoder(device, fragmentArgs)))
return nullptr;
- if (computeArgsArray && !(compute = tryCreateMtlArgumentEncoder(device, computeArgsArray)))
+ if (computeArgs && !(compute = tryCreateMtlArgumentEncoder(device, computeArgs)))
return nullptr;
return adoptRef(new GPUBindGroupLayout(WTFMove(bindingsMap), WTFMove(vertex), WTFMove(fragment), WTFMove(compute)));
