Attachment 373056 Details for Bug 199093 – WIP

[patch] WIP

b-backup.diff (text/plain), 11.05 KB, created by Saam Barati on 2019-06-27 15:12:35 PDT

(hide)

Description:

Filename:

MIME Type:

Creator: Saam Barati

Created: 2019-06-27 15:12:35 PDT

Size: 11.05 KB

patch

obsolete

>Index: Source/WebCore/Modules/webgpu/WHLSL/WHLSLPrepare.cpp
>===================================================================
>--- Source/WebCore/Modules/webgpu/WHLSL/WHLSLPrepare.cpp	(revision 246889)
>+++ Source/WebCore/Modules/webgpu/WHLSL/WHLSLPrepare.cpp	(working copy)
>@@ -60,7 +60,7 @@ namespace WHLSL {
> 
> static constexpr bool dumpASTBeforeEachPass = false;
> static constexpr bool dumpASTAfterParsing = false;
>-static constexpr bool dumpASTAtEnd = false;
>+static constexpr bool dumpASTAtEnd = true;
> static constexpr bool alwaysDumpPassFailures = false;
> static constexpr bool dumpPassFailure = dumpASTBeforeEachPass || dumpASTAfterParsing || dumpASTAtEnd || alwaysDumpPassFailures;
> 
>Index: Source/WebCore/Modules/webgpu/WHLSL/WHLSLStandardLibrary.txt
>===================================================================
>--- Source/WebCore/Modules/webgpu/WHLSL/WHLSLStandardLibrary.txt	(revision 246889)
>+++ Source/WebCore/Modules/webgpu/WHLSL/WHLSLStandardLibrary.txt	(working copy)
>@@ -744,6 +744,12 @@ operator uint2(uint x, uint y) {
>     return result;
> }
> 
>+native bool operator==(uchar, uchar);
>+native bool operator==(ushort, ushort);
>+native bool operator==(char, char);
>+native bool operator==(short, short);
>+native bool operator==(half, half);
>+
> native int operator+(int, int);
> native int operator-(int, int);
> native int operator*(int, int);
>@@ -1172,4 +1178,52 @@ native void GetDimensions(Texture2D<floa
> native void GetDimensions(Texture2D<float4>, uint MipLevel, threadgroup uint* Width, threadgroup uint* Height, device uint* NumberOfLevels);
> native void GetDimensions(Texture2D<float4>, uint MipLevel, threadgroup uint* Width, threadgroup uint* Height, threadgroup uint* NumberOfLevels);
> 
>+
>+native bool operator==(bool, bool);
>+native bool operator&&(bool, bool);
>+native bool operator||(bool, bool);
>+bool operator&(bool a, bool b) {
>+    return a && b;
>+}
>+bool operator|(bool a, bool b) {
>+    return a || b;
>+}
>+bool operator^(bool a, bool b) {
>+    if (a)
>+        return !b;
>+    return b;
>+}
>+bool operator~(bool value) {
>+    return !value;
>+}
>+
>+operator bool(uchar x) {
>+    return x != 0;
>+}
>+operator bool(ushort x) {
>+    return x != 0;
>+}
>+operator bool(uint x) {
>+    return x != 0;
>+}
>+operator bool(char x) {
>+    return x != 0;
>+}
>+operator bool(short x) {
>+    return x != 0;
>+}
>+operator bool(int x) {
>+    return x != 0;
>+}
>+operator bool(half x) {
>+    return x != 0;
>+}
>+operator bool(float x) {
>+    return x != 0;
>+}
>+
>+operator int(bool x) {
>+    return x ? 1 : 0;
>+}
>+
> // FIXME: https://bugs.webkit.org/show_bug.cgi?id=192890 Insert the rest of the standard library once the parser is fast enough
>Index: LayoutTests/webgpu/whlsl-bitwise-bool-ops.html
>===================================================================
>--- LayoutTests/webgpu/whlsl-bitwise-bool-ops.html	(nonexistent)
>+++ LayoutTests/webgpu/whlsl-bitwise-bool-ops.html	(working copy)
>@@ -0,0 +1,155 @@
>+<!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;
>+
>+const whlslTests = {};
>+
>+whlslTests.boolBitAnd = () => {
>+    const source = `
>+        bool foo(bool a, bool b)
>+        {
>+            return a & b;
>+        }
>+    `;
>+
>+    webGPUPromiseTest(async () => {
>+        try {
>+            {
>+                let result = await callBoolFunction(source, "foo", [makeBool(true), makeBool(true)]);
>+                assert_equals(result, true, "Test returned expected value.");
>+            }
>+
>+            {
>+                let result = await callBoolFunction(source, "foo", [makeBool(false), makeBool(false)]);
>+                assert_equals(result, false, "Test returned expected value.");
>+            }
>+
>+            {
>+                let result = await callBoolFunction(source, "foo", [makeBool(true), makeBool(false)]);
>+                assert_equals(result, false, "Test returned expected value.");
>+            }
>+
>+            {
>+                let result = await callBoolFunction(source, "foo", [makeBool(false), makeBool(true)]);
>+                assert_equals(result, false, "Test returned expected value.");
>+            }
>+        } catch(e) {
>+            if (!(e instanceof WebGPUUnsupportedError))
>+                throw e;
>+        }
>+    }, "Bool bit and");
>+};
>+
>+whlslTests.boolBitOr = () => {
>+    const source = `
>+        bool foo(bool a, bool b)
>+        {
>+            return a | b;
>+        }
>+    `;
>+
>+    webGPUPromiseTest(async () => {
>+        try {
>+            {
>+                let result = await callBoolFunction(source, "foo", [makeBool(true), makeBool(true)]);
>+                assert_equals(result, true, "Test returned expected value.");
>+            }
>+
>+            {
>+                let result = await callBoolFunction(source, "foo", [makeBool(false), makeBool(false)]);
>+                assert_equals(result, false, "Test returned expected value.");
>+            }
>+
>+            {
>+                let result = await callBoolFunction(source, "foo", [makeBool(true), makeBool(false)]);
>+                assert_equals(result, true, "Test returned expected value.");
>+            }
>+
>+            {
>+                let result = await callBoolFunction(source, "foo", [makeBool(false), makeBool(true)]);
>+                assert_equals(result, true, "Test returned expected value.");
>+            }
>+        } catch(e) {
>+            if (!(e instanceof WebGPUUnsupportedError))
>+                throw e;
>+        }
>+    }, "Bool bit or");
>+};
>+
>+whlslTests.boolBitNot = () => {
>+    const source = `
>+        bool foo(bool a)
>+        {
>+            return ~a;
>+        }
>+    `;
>+
>+    webGPUPromiseTest(async () => {
>+        try {
>+            {
>+                let result = await callBoolFunction(source, "foo", [makeBool(true)]);
>+                assert_equals(result, false, "Test returned expected value.");
>+            }
>+
>+            {
>+                let result = await callBoolFunction(source, "foo", [makeBool(false)]);
>+                assert_equals(result, true, "Test returned expected value.");
>+            }
>+        } catch(e) {
>+            if (!(e instanceof WebGPUUnsupportedError))
>+                throw e;
>+        }
>+    }, "Bool bit not");
>+};
>+
>+function runTests(obj) {
>+    window.addEventListener("load", () => {
>+        try {
>+            for (const name in obj) {
>+                if (!name.startsWith("_"))
>+                    obj[name]();
>+            }
>+        } catch (e) {
>+            if (window.testRunner)
>+                testRunner.notifyDone();
>+            
>+            throw e;
>+        }
>+    });
>+}
>+
>+runTests(whlslTests);
>+
>+const checkBools = (msg = "Return an expected bool value.", body, argValues = [], expected = true) => {
>+    // FIXME (https://webkit.org/b/199093): Bool[] functions don't compile, so no-op for now.
>+    return;
>+
>+    const [src, name, values] = appendScalarFunctionToSource("", "bool", body, argValues); 
>+
>+    webGPUPromiseTest(async () => {
>+        return callBoolFunction(src, name, values).then(result => {
>+            assert_equals(result, expected, "Test returned expected value.");
>+        }, e => {
>+            if (!(e instanceof WebGPUUnsupportedError))
>+                throw e;
>+        });
>+    }, msg);
>+};
>+
>+const webGPUPromiseTest = (testFunc, msg) => {
>+    promise_test(async () => { 
>+        return testFunc().catch(e => {
>+        if (!(e instanceof WebGPUUnsupportedError))
>+            throw e;
>+        });
>+    }, msg);
>+}
>+</script>
>+</html>
>Index: LayoutTests/webgpu/whlsl-compute.html
>===================================================================
>--- LayoutTests/webgpu/whlsl-compute.html	(revision 246889)
>+++ LayoutTests/webgpu/whlsl-compute.html	(working copy)
>@@ -7,6 +7,11 @@
> <body>
> <script>
> const shaderSource = `
>+bool foo(int a, int b)
>+{
>+    return bool(a) & bool(b);
>+}
>+
> [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;
>Index: LayoutTests/webgpu/js/whlsl-test-harness.js
>===================================================================
>--- LayoutTests/webgpu/js/whlsl-test-harness.js	(revision 246889)
>+++ LayoutTests/webgpu/js/whlsl-test-harness.js	(working copy)
>@@ -26,7 +26,7 @@ function convertTypeToArrayType(type)
> {
>     switch(type) {
>         case Types.BOOL:
>-            return Uint8Array;
>+            return Int32Array;
>         case Types.INT:
>             return Int32Array;
>         case Types.UCHAR:
>@@ -61,6 +61,27 @@ function convertTypeToWHLSLType(type)
>     }
> }
> 
>+function whlslArgumentType(type)
>+{
>+    if (type === Types.BOOL)
>+        return "int";
>+    return convertTypeToWHLSLType(type);
>+}
>+
>+function convertToOutputType(code, type)
>+{
>+    if (type !== Types.BOOL)
>+        return code;
>+    return `int(${code})`;
>+}
>+
>+function convertToInputType(code, type)
>+{
>+    if (type !== Types.BOOL)
>+        return code;
>+    return `bool(${code})`;
>+}
>+
> /* Harness Classes */
> 
> class WebGPUUnsupportedError extends Error {
>@@ -196,12 +217,14 @@ using namespace metal;
> 
>         let entryPointCode;
>         if (this._isWHLSL) {
>-            argsDeclarations.unshift(`device ${convertTypeToWHLSLType(type)}[] result : register(u0)`);
>+            argsDeclarations.unshift(`device ${whlslArgumentType(type)}[] result : register(u0)`);
>+            let callCode = `${name}(${functionCallArgs.join(", ")})`;
>+            callCode = convertToOutputType(callCode, type);
>             entryPointCode = `
> [numthreads(1, 1, 1)]
> compute void _compute_main(${argsDeclarations.join(", ")})
> {
>-    result[0] = ${name}(${functionCallArgs.join(", ")});
>+    result[0] = ${callCode};
> }
> `;
>         } else {
>@@ -218,6 +241,7 @@ kernel void _compute_main(device _comput
> `;
>         }
>         const code = this._shaderHeader + functions + entryPointCode;
>+        console.log(code);
>         this._callFunction(code, argsLayouts, argsResourceBindings);
>     
>         try {
>@@ -301,8 +325,8 @@ kernel void _compute_main(device _comput
>         for (let i = 1; i <= args.length; ++i) {
>             const arg = args[i - 1];
>             if (this._isWHLSL) {
>-                argsDeclarations.push(`device ${convertTypeToWHLSLType(arg.type)}[] arg${i} : register(u${i})`);
>-                functionCallArgs.push(`arg${i}` + (arg.isBuffer ? "" : "[0]"));
>+                argsDeclarations.push(`device ${whlslArgumentType(arg.type)}[] arg${i} : register(u${i})`);
>+                functionCallArgs.push(convertToInputType(`arg${i}` + (arg.isBuffer ? "" : "[0]"), arg.type));
>             } else {
>                 argsDeclarations.push(`device ${convertTypeToWHLSLType(arg.type)}* arg${i} [[id(${i})]];`);
>                 functionCallArgs.push((arg.isBuffer ? "" : "*") + `args.arg${i}`);
>@@ -456,4 +480,4 @@ async function callFloat4Function(functi
> function callVoidFunction(functions, name, args)
> {
>     harness.callVoidFunction(functions, name, args);
>-}
>\ No newline at end of file
>+}

Actions: View | Formatted Diff | Diff

Attachments on bug 199093: 373056 | 373248