Attachment 372666 Details for Bug 199028 – Patch

[patch] Patch

bug-199028-20190621190533.patch (text/plain), 41.83 KB, created by Justin Fan on 2019-06-21 19:05:34 PDT

(hide)

Description:

Filename:

MIME Type:

Creator: Justin Fan

Created: 2019-06-21 19:05:34 PDT

Size: 41.83 KB

patch

obsolete

>Subversion Revision: 246704
>diff --git a/LayoutTests/ChangeLog b/LayoutTests/ChangeLog
>index 96ab87cc52fdfdd1696136f8939488dff0408ee5..2d475b6f5c3d2b8acc57fb91ff70b90cac29032a 100644
>--- a/LayoutTests/ChangeLog
>+++ b/LayoutTests/ChangeLog
>@@ -1,3 +1,40 @@
>+2019-06-21  Justin Fan  <justin_fan@apple.com>
>+
>+        [WHLSL] Make whlsl-test-harness actually generate WHLSL shaders by default
>+        https://bugs.webkit.org/show_bug.cgi?id=199028
>+
>+        Reviewed by NOBODY (OOPS!).
>+
>+        whlsl-test-harness.js now generates WHLSL shaders and invokes WebKit's WHLSL compiler.
>+        MSL mode remains to facilitate further harness tesing.
>+        In addition, if WebGPU is not supported, synchronous Harness methods do nothing.
>+        Asynchronous methods will throw a WebGPUUnsupportedError that "rejects" the returned Promise.
>+
>+        * TestExpectations:
>+        * platform/mac/TestExpectations:
>+        * webgpu/js/whlsl-test-harness.js:
>+        (WebGPUUnsupportedError): Layout tests should catch these to fail gracefully if WebGPU is not supported.
>+        (Data):
>+        (Data.prototype.async.getArrayBuffer):
>+        (Data.prototype.get isBuffer): Renamed from isPointer.
>+        (Harness):
>+        (Harness.prototype.async.requestDevice): Can be used to re-acquire a GPUDevice.
>+        (Harness.prototype.set isWHLSL): Determines whether harness will generate WHLSL or MSL shaders.
>+        (Harness.prototype.async.callTypedFunction):
>+        (Harness.prototype.callVoidFunction):
>+        (Harness.prototype.get device):
>+        (Harness.prototype._clearResults):
>+        (Harness.prototype._setUpArguments):
>+        (Harness.prototype._callFunction):
>+        (Data.prototype.get isPointer): Deleted.
>+        (Harness.prototype._initialize): Deleted.
>+        (Harness.prototype.async.callVoidFunction): Deleted.
>+        (harness._initialize.async): Deleted.
>+        * webgpu/msl-harness-test-expected.txt: Renamed from LayoutTests/webgpu/whlsl-harness-test-expected.txt.
>+        * webgpu/msl-harness-test.html: Copied from LayoutTests/webgpu/whlsl-harness-test.html.
>+        * webgpu/whlsl-test-harness-test-expected.html: Added.
>+        * webgpu/whlsl-test-harness-test.html: Renamed from LayoutTests/webgpu/whlsl-harness-test.html.
>+
> 2019-06-21  Saam Barati  <sbarati@apple.com>
> 
>         [WHLSL] Code that accesses an undefined variable crashes
>diff --git a/LayoutTests/TestExpectations b/LayoutTests/TestExpectations
>index e73d995fae93fef508a2138ad83e1d8dace3c7dd..3df46cfc527b9d492d4923c04091ca3b02ed719c 100644
>--- a/LayoutTests/TestExpectations
>+++ b/LayoutTests/TestExpectations
>@@ -3438,3 +3438,5 @@ imported/w3c/web-platform-tests/websockets/Secure-Send-unpaired-surrogates.any.w
> 
> # iOS only
> fast/dom/linkify-phone-numbers.html [ ImageOnlyFailure ]
>+
>+webkit.org/b/199028 webgpu/whlsl-test-harness-test.html [ Slow ]
>diff --git a/LayoutTests/platform/mac/TestExpectations b/LayoutTests/platform/mac/TestExpectations
>index 1164b451d0b6cc6af251592ca2abfc8463f17dbb..0fa432dd2e2e41300b8d994c5c9cbe0792676a03 100644
>--- a/LayoutTests/platform/mac/TestExpectations
>+++ b/LayoutTests/platform/mac/TestExpectations
>@@ -1787,7 +1787,8 @@ webkit.org/b/192956 [ Sierra ] webgpu [ Skip ]
> webkit.org/b/192956 [ Sierra ] inspector/canvas/create-context-webgpu.html [ Skip ]
> webkit.org/b/192956 [ Sierra ] inspector/canvas/resolveCanvasContext-webgpu.html [ Skip ]
> 
>-webkit.org/b/199076 [ HighSierra ] webgpu/whlsl-harness-test.html [ Skip ]
>+webkit.org/b/199076 [ HighSierra ] webgpu/whlsl-test-harness-test.html [ Skip ]
>+webkit.org/b/199028 [ HighSierra ] webgpu/msl-harness-test.html [ Skip ]
> 
> webkit.org/b/189680 platform/mac/media/audio-session-category-video-paused.html [ Pass Timeout ]
> 
>diff --git a/LayoutTests/webgpu/js/whlsl-test-harness.js b/LayoutTests/webgpu/js/whlsl-test-harness.js
>index 2c9213044a152471287e3ef55733b6ea9c368453..1f1aa5e9f9f0805db06ee97a56c6c48826d83775 100644
>--- a/LayoutTests/webgpu/js/whlsl-test-harness.js
>+++ b/LayoutTests/webgpu/js/whlsl-test-harness.js
>@@ -63,28 +63,37 @@ function convertTypeToWHLSLType(type)
> 
> /* Harness Classes */
> 
>+class WebGPUUnsupportedError extends Error {
>+    constructor()
>+    {
>+        super("No GPUDevice detected!");
>+    }
>+};
>+
> 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)
>+    constructor(harness, type, values, isBuffer = false)
>     {
>-        // One or more scalars in an array can be accessed through a pointer to buffer.
>+        if (harness.device === undefined)
>+            return;
>+        // One or more scalars in an array can be accessed through an array reference.
>         // 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;
>+            this._isBuffer = isVectorType(type) ? isBuffer : true;
>         else {
>-            this._isPointer = false;
>+            this._isBuffer = false;
>             values = [values];
>         }
> 
>         this._type = type;
>         this._byteLength = (convertTypeToArrayType(type)).BYTES_PER_ELEMENT * values.length;
> 
>-        const [buffer, arrayBuffer] = harness._device.createBufferMapped({
>+        const [buffer, arrayBuffer] = harness.device.createBufferMapped({
>             size: this._byteLength,
>             usage: GPUBufferUsage.STORAGE | GPUBufferUsage.MAP_READ
>         });
>@@ -101,6 +110,9 @@ class Data {
>      */
>     async getArrayBuffer()
>     {
>+        if (harness.device === undefined)
>+            throw new WebGPUUnsupportedError();
>+
>         let result;
>         try {
>             result = await this._buffer.mapReadAsync();
>@@ -112,22 +124,37 @@ class Data {
>     }
> 
>     get type() { return this._type; }
>-    get isPointer() { return this._isPointer; }
>+    get isBuffer() { return this._isBuffer; }
>     get buffer() { return this._buffer; }
>     get byteLength() { return this._byteLength; }
> }
> 
> class Harness {
>-    constructor()
>+    constructor ()
>     {
>-        this._shaderHeader = `#include <metal_stdlib>
>-        using namespace metal;
>-        `;
>+        this.isWHLSL = true;
>     }
> 
>-    _initialize(callback)
>+    async requestDevice()
>     {
>-        callback.bind(this)();
>+        try {
>+            const adapter = await navigator.gpu.requestAdapter();
>+            this._device = await adapter.requestDevice();
>+        } catch {
>+            // WebGPU is not supported.
>+            // FIXME: Add support for GPUAdapterRequestOptions and GPUDeviceDescriptor,
>+            // and differentiate between descriptor validation errors and no WebGPU support.
>+        }
>+    }
>+
>+    // Sets whether Harness generates WHLSL or MSL shaders.
>+    set isWHLSL(value)
>+    {
>+        this._isWHLSL = value;
>+        this._shaderHeader = value ? "" : `
>+#include <metal_stdlib>
>+using namespace metal;
>+        `;
>     }
> 
>     /**
>@@ -140,12 +167,17 @@ class Harness {
>      */
>     async callTypedFunction(type, functions, name, args)
>     {   
>-        const [argsLayouts, argsResourceBindings, argsStructCode, functionCallArgs] = this._setUpArguments(args);
>+        if (this._device === undefined)
>+            throw new WebGPUUnsupportedError();
> 
>-        if (!this._resultBuffer) {
>-            this._resultBuffer = this._device.createBuffer({ 
>+        const [argsLayouts, argsResourceBindings, argsDeclarations, functionCallArgs] = this._setUpArguments(args);
>+
>+        if (this._resultBuffer) {
>+            this._clearResults()
>+        } else {
>+            this._resultBuffer = this.device.createBuffer({ 
>                 size: Types.MAX_SIZE, 
>-                usage: GPUBufferUsage.STORAGE | GPUBufferUsage.MAP_READ 
>+                usage: GPUBufferUsage.STORAGE | GPUBufferUsage.MAP_READ | GPUBufferUsage.TRANSFER_DST
>             });
>         }
> 
>@@ -162,17 +194,30 @@ class Harness {
>             }
>         });
> 
>-        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(", ")});
>+        let entryPointCode;
>+        if (this._isWHLSL) {
>+            argsDeclarations.unshift(`device ${convertTypeToWHLSLType(type)}[] result : register(u0)`);
>+            entryPointCode = `
>+[numthreads(1, 1, 1)]
>+compute void _compute_main(${argsDeclarations.join(", ")})
>+{
>+    result[0] = ${name}(${functionCallArgs.join(", ")});
>+}
>+`;
>+        } else {
>+            argsDeclarations.unshift(`device ${convertTypeToWHLSLType(type)}* result [[id(0)]];`);
>+            entryPointCode = `
>+struct _compute_args {
>+    ${argsDeclarations.join("\n")}
>+};
>+
>+kernel void _compute_main(device _compute_args& args [[buffer(0)]]) 
>+{
>+    *args.result = ${name}(${functionCallArgs.join(", ")});
>+}
>+`;
>         }
>-        `;
>-    
>+        const code = this._shaderHeader + functions + entryPointCode;
>         this._callFunction(code, argsLayouts, argsResourceBindings);
>     
>         try {
>@@ -192,23 +237,52 @@ class Harness {
>      * @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)
>+    callVoidFunction(functions, name, args)
>     {
>-        const [argsLayouts, argsResourceBindings, argsStructCode, functionCallArgs] = this._setUpArguments(args);
>+        if (this._device === undefined)
>+            return;
> 
>-        const code = this._shaderHeader + functions + `
>-        struct _compute_args {
>-            ${argsStructCode}};
>+        const [argsLayouts, argsResourceBindings, argsDeclarations, functionCallArgs] = this._setUpArguments(args);
> 
>-        kernel void _compute_main(device _compute_args& args [[buffer(0)]])
>-        {
>-            ${name}(${functionCallArgs.join(", ")});
>+        let entryPointCode;
>+        if (this._isWHLSL) {
>+            entryPointCode = `
>+[numthreads(1, 1, 1)]
>+compute void _compute_main(${argsDeclarations.join(", ")})
>+{
>+    ${name}(${functionCallArgs.join(", ")});
>+}`;
>+        } else {
>+            entryPointCode = `
>+struct _compute_args {
>+    ${argsDeclarations.join("\n")}
>+};
>+
>+kernel void _compute_main(device _compute_args& args [[buffer(0)]])
>+{
>+    ${name}(${functionCallArgs.join(", ")});
>+}
>+`;
>         }
>-        `;
>-
>+        const code = this._shaderHeader + functions + entryPointCode;
>         this._callFunction(code, argsLayouts, argsResourceBindings);
>     }
> 
>+    get device() { return this._device; }
>+
>+    _clearResults()
>+    {
>+        if (!this._clearBuffer) {
>+            this._clearBuffer = this._device.createBuffer({ 
>+                size: Types.MAX_SIZE, 
>+                usage: GPUBufferUsage.TRANSFER_SRC
>+            });
>+        }
>+        const commandEncoder = this._device.createCommandEncoder();
>+        commandEncoder.copyBufferToBuffer(this._clearBuffer, 0, this._resultBuffer, 0, Types.MAX_SIZE);
>+        this._device.getQueue().submit([commandEncoder.finish()]);
>+    }
>+
>     _setUpArguments(args)
>     {
>         if (!Array.isArray(args)) {
>@@ -220,17 +294,20 @@ class Harness {
> 
>         // FIXME: Replace with WHLSL.
>         // Expand bind group structure to represent any arguments.
>-        let argsStructCode = "";
>+        let argsDeclarations = [];
>         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}`);
>+            if (this._isWHLSL) {
>+                argsDeclarations.push(`device ${convertTypeToWHLSLType(arg.type)}[] arg${i} : register(u${i})`);
>+                functionCallArgs.push(`arg${i}` + (arg.isBuffer ? "" : "[0]"));
>+            } else {
>+                argsDeclarations.push(`device ${convertTypeToWHLSLType(arg.type)}* arg${i} [[id(${i})]];`);
>+                functionCallArgs.push((arg.isBuffer ? "" : "*") + `args.arg${i}`);
>+            }
>             argsLayouts.push({
>                 binding: i,
>                 visibility: GPUShaderStageBit.COMPUTE,
>@@ -245,29 +322,33 @@ class Harness {
>             });
>         }
> 
>-        return [argsLayouts, argsResourceBindings, argsStructCode, functionCallArgs];
>+        return [argsLayouts, argsResourceBindings, argsDeclarations, functionCallArgs];
>     }
> 
>     _callFunction(code, argsLayouts, argsResourceBindings)
>     {
>-        const shaders = this._device.createShaderModule({ code: code });
>+        const shaders = this._device.createShaderModule({ code: code, isWHLSL: this._isWHLSL });
>+
>+        const bindGroupLayout = this._device.createBindGroupLayout({
>+            bindings: argsLayouts
>+        });
>+
>+        const pipelineLayout = this._device.createPipelineLayout({ bindGroupLayouts: [bindGroupLayout] });
>+
>+        const bindGroup = this._device.createBindGroup({
>+            layout: bindGroupLayout,
>+            bindings: argsResourceBindings
>+        });
>+
>         // FIXME: Compile errors should be caught and reported here.
>         const pipeline = this._device.createComputePipeline({
>+            layout: pipelineLayout,
>             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);
>@@ -282,14 +363,7 @@ class Harness {
> /* 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!");
>-    }
>-});
>+harness.requestDevice();
> 
> /* Global Helper Functions */
> 
>diff --git a/LayoutTests/webgpu/msl-harness-test-expected.txt b/LayoutTests/webgpu/msl-harness-test-expected.txt
>new file mode 100644
>index 0000000000000000000000000000000000000000..e87095632c61eae31fb81dd47d0507d2b3d8bdc0
>--- /dev/null
>+++ b/LayoutTests/webgpu/msl-harness-test-expected.txt
>@@ -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*. 
>+
>diff --git a/LayoutTests/webgpu/msl-harness-test.html b/LayoutTests/webgpu/msl-harness-test.html
>new file mode 100644
>index 0000000000000000000000000000000000000000..8d43fce9eb7097a94c2c28829fc429033ec4bf53
>--- /dev/null
>+++ b/LayoutTests/webgpu/msl-harness-test.html
>@@ -0,0 +1,218 @@
>+<!DOCTYPE html>
>+<html>
>+<meta charset=utf-8>
>+<title>Test the WHLSL test harness.</title>
>+<script src="js/whlsl-test-harness.js"></script>
>+<script src="../resources/testharness.js"></script>
>+<script src="../resources/testharnessreport.js"></script>
>+<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 mslTests = {};
>+
>+mslTests.literals = () => {
>+    checkBools("Return a literal of type bool.", "return true;");
>+    checkFloat4s("return float4(0, 1, 2, 3);");
>+    checkNumericScalars("return 42;", [], 42);
>+};
>+
>+mslTests.singleArgument = () => {
>+    checkBools("Upload and return a bool value.", "return in0;", [true]);
>+    checkFloat4s("return in0.wzyx;", [[3, 2, 1, 0]]);
>+    checkNumericScalars("return in0;", [42], 42);
>+};
>+
>+mslTests.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);
>+};
>+
>+mslTests.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);
>+};
>+
>+mslTests.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);
>+};
>+
>+mslTests.multipleArgumentTypes = () => {
>+    const src = `float 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]);
>+    webGPUPromiseTest(() => {
>+        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.");
>+};
>+
>+mslTests.bufferStores = () => {
>+    let src = `void test(device float4* out) {
>+        *out = float4(0, 1, 2, 3);
>+    }`;
>+    const float4Out = makeFloat4([[0, 0, 0, 0]]);
>+    callVoidFunction(src, "test", float4Out);
>+
>+    webGPUPromiseTest(() => {
>+        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 = `void 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]);
>+
>+    webGPUPromiseTest(() => {
>+        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", () => {
>+    harness.isWHLSL = false;
>+    for (const name in mslTests) {
>+        mslTests[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];
>+        webGPUPromiseTest(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} }
>+    `;
>+
>+    webGPUPromiseTest(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} }
>+    `;
>+
>+    webGPUPromiseTest(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.");
>+}
>+
>+const webGPUPromiseTest = (testFunc, msg) => {
>+    promise_test(async () => { 
>+        return testFunc().catch(e => {
>+        if (!(e instanceof WebGPUUnsupportedError))
>+            throw e;
>+        });
>+    }, msg);
>+}
>+</script>
>+</html>
>\ No newline at end of file
>diff --git a/LayoutTests/webgpu/whlsl-harness-test-expected.txt b/LayoutTests/webgpu/whlsl-harness-test-expected.txt
>deleted file mode 100644
>index e87095632c61eae31fb81dd47d0507d2b3d8bdc0..0000000000000000000000000000000000000000
>--- a/LayoutTests/webgpu/whlsl-harness-test-expected.txt
>+++ /dev/null
>@@ -1,35 +0,0 @@
>-
>-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*. 
>-
>diff --git a/LayoutTests/webgpu/whlsl-harness-test.html b/LayoutTests/webgpu/whlsl-harness-test.html
>deleted file mode 100644
>index c00081d7834f7bc742fd37c20802aead0be79e85..0000000000000000000000000000000000000000
>--- a/LayoutTests/webgpu/whlsl-harness-test.html
>+++ /dev/null
>@@ -1,208 +0,0 @@
>-<!DOCTYPE html>
>-<html>
>-<meta charset=utf-8>
>-<title>Test the WHLSL test harness.</title>
>-<script src="js/whlsl-test-harness.js"></script>
>-<script src="../resources/testharness.js"></script>
>-<script src="../resources/testharnessreport.js"></script>
>-<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 = `float 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 = `void 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 = `void 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
>diff --git a/LayoutTests/webgpu/whlsl-test-harness-test.html b/LayoutTests/webgpu/whlsl-test-harness-test.html
>new file mode 100644
>index 0000000000000000000000000000000000000000..e368c6b628013f676a87f06c36b73a4bbcd3b05c
>--- /dev/null
>+++ b/LayoutTests/webgpu/whlsl-test-harness-test.html
>@@ -0,0 +1,242 @@
>+<!DOCTYPE html>
>+<html>
>+<meta charset=utf-8>
>+<meta name="timeout" content="long">
>+<title>Test the WHLSL test harness.</title>
>+<script src="js/whlsl-test-harness.js"></script>
>+<script src="../resources/testharness.js"></script>
>+<script src="../resources/testharnessreport.js"></script>
>+<script>
>+const epsilon = 0.0001;
>+
>+// FIXME: Add "uchar" back in when operator+(uchar, uchar) is available.
>+const numericScalarTypes = ["int", "uint", "float"];
>+
>+const numericScalarFuncs = {
>+    "int": callIntFunction,
>+    "uchar": callUcharFunction,
>+    "uint": callUintFunction,
>+    "float": callFloatFunction
>+};
>+
>+const scalarArgMakers = {
>+    "int": makeInt,
>+    "uchar": makeUchar,
>+    "uint": makeUint,
>+    "float": makeFloat
>+};
>+
>+let whlslTests = {};
>+
>+whlslTests.literals = () => {
>+    checkBools("Return a literal of type bool.", "return true;");
>+    checkFloat4s("return float4(0, 1, 2, 3);");
>+    checkNumericScalars("return 42;", [], 42);
>+};
>+
>+whlslTests.singleArgument = () => {
>+    checkBools("Upload and return a bool value.", "return in0;", [true]);
>+    checkFloat4s("return in0;", [[0, 1, 2, 3]]);
>+    checkNumericScalars("return in0;", [42], 42);
>+};
>+
>+whlslTests.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]);
>+
>+    let body = "return float4(in0.x, in1.y, in2.z, in3.w);"
>+    let args = [];
>+    for (let i = 0; i < 4; ++i)
>+        args.push([0, 1, 2, 3]);
>+    checkFloat4s(body, args, [0, 1, 2, 3]);
>+
>+    body = `return in0 + in1 + in2 + in3 + in4 + in5 + in6 + in7;`;
>+    checkNumericScalars(body, [0, 1, 2, 3, 4, 5, 6, 7], 28);
>+};
>+
>+whlslTests.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);
>+};
>+
>+whlslTests.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]]);
>+
>+    let body = `
>+    float x = in0[0].x + in0[1].x + in0[2].x + in1.x + in2[0].x;
>+    float y = in0[0].y + in0[1].y + in0[2].y + in1.y + in2[0].y;
>+    float z = in0[0].z + in0[1].z + in0[2].z + in1.z + in2[0].z;
>+    float w = in0[0].w + in0[1].w + in0[2].w + in1.w + in2[0].w;
>+
>+    return float4(x, y, z, w);`;
>+    const vector = [0, 1, 2, 3];
>+    checkFloat4s(body, [[vector, vector, vector], vector, [vector]], [0, 5, 10, 15]);
>+
>+    body = `return in0[0] + in0[1] + in0[2] + in1 + in2[0];`;
>+    checkNumericScalars(body, [[0, 1, 2], 3, [4]], 10);
>+};
>+
>+// This test requires a whole lot of missing operator+ and operator&& so skip for now.
>+whlslTests._disabled_multipleArgumentTypes = () => {
>+    const src = `float 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]);
>+    webGPUPromiseTest(() => {
>+        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.");
>+};
>+
>+whlslTests.bufferStores = () => {
>+    let src = `void test(device float4[] out) {
>+        out[0] = float4(0, 1, 2, 3);
>+    }`;
>+    const float4Out = makeFloat4([[0, 0, 0, 0]]);
>+    callVoidFunction(src, "test", float4Out);
>+
>+    webGPUPromiseTest(() => {
>+        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 = `void test(device int[] in, device int[] out) {
>+        for (uint i = 0; i < 5; i = i + 1)
>+            out[i] = in[i];
>+    }`;
>+    const array = [0, 1, 2, 3, 4];
>+    const input = makeInt(array);
>+    const output = makeInt([0, 0, 0, 0, 0]);
>+    callVoidFunction(src, "test", [input, output]);
>+
>+    webGPUPromiseTest(() => {
>+        return output.getArrayBuffer().then(arrayBuffer => {
>+            const result = new Uint32Array(arrayBuffer);
>+            assert_array_equals(array, result, "Test stored expected values.");
>+        });
>+    }, "Upload a int[] and store into a int[].");
>+};
>+
>+window.addEventListener("load", () => {
>+    try {
>+        for (const name in whlslTests) {
>+            if (!name.startsWith("_"))
>+                whlslTests[name]();
>+        }
>+    } catch {
>+        if (window.testRunner)
>+            testRunner.notifyDone();
>+    }
>+});
>+
>+/* 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 isBuffer = Array.isArray(argValues[i]);
>+            inputArgs.push(`${isBuffer ? "device " : ""}${type}${isBuffer ? "[]" : ""} 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];
>+        webGPUPromiseTest(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) => {
>+    // Bool functions don't compile due to bug https://bugs.webkit.org/show_bug.cgi?id=199093, so no-op for now.
>+    return;
>+
>+    let src = "";
>+    let inputArgs = [];
>+    let values = [];
>+    for (let i = 0; i < argValues.length; ++i) {
>+        const isBuffer = Array.isArray(argValues[i]);
>+        inputArgs.push(`${isBuffer ? "device " : ""}bool${isBuffer ? "[]" : ""} in${i}`);
>+        values.push(makeBool(argValues[i]));
>+    }
>+
>+    src += `bool boolTest(${inputArgs.join(", ")}) { ${body} }
>+    `;
>+
>+    webGPUPromiseTest(async () => {
>+        return callBoolFunction(src, "boolTest", values).then(result => {
>+            assert_equals(result, expected, "Test returned expected value.");
>+        }, e => {
>+            if (!(e instanceof WebGPUUnsupportedError))
>+                throw e;
>+        });
>+    }, 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 isBuffer = argValues[i].length === 1 || totalLength > 4;
>+        inputArgs.push(`${isBuffer ? "device " : ""}float4${isBuffer ? "[]" : ""} in${i}`);
>+        values.push(makeFloat4(argValues[i]));
>+    }
>+
>+    src += `float4 float4Test(${inputArgs.join(", ")}) { ${body} }
>+    `;
>+
>+    webGPUPromiseTest(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.");
>+}
>+
>+const webGPUPromiseTest = (testFunc, msg) => {
>+    promise_test(async () => { 
>+        return testFunc().catch(e => {
>+        if (!(e instanceof WebGPUUnsupportedError))
>+            throw e;
>+        });
>+    }, msg);
>+}
>+</script>
>+</html>
>\ No newline at end of file

Actions: View | Formatted Diff | Diff

Attachments on bug 199028: 372666 | 372670 | 372884 | 372885