Attachment 372420 Details for Bug 198978 – Patch

[patch] Patch

bug-198978-20190618194310.patch (text/plain), 30.34 KB, created by Justin Fan on 2019-06-18 19:43:11 PDT

(hide)

Description:

Filename:

MIME Type:

Creator: Justin Fan

Created: 2019-06-18 19:43:11 PDT

Size: 30.34 KB

patch

obsolete

>Subversion Revision: 246578
>diff --git a/Source/WebCore/ChangeLog b/Source/WebCore/ChangeLog
>index dcfefd4e27a6e8c25fc45fa865187f8ac4183dc9..5eb8347871c54bdeb3622cffe2402e7e875f26bf 100644
>--- a/Source/WebCore/ChangeLog
>+++ b/Source/WebCore/ChangeLog
>@@ -1,3 +1,21 @@
>+2019-06-18  Justin Fan  <justin_fan@apple.com>
>+
>+        [WHLSL] Create a shading language test harness
>+        https://bugs.webkit.org/show_bug.cgi?id=198978
>+
>+        Reviewed by NOBODY (OOPS!).
>+
>+        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-18  Yusuke Suzuki  <ysuzuki@apple.com>
> 
>         [JSC] JSLock should be WebThread aware
>diff --git a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp
>index ec45da13f86e34248eade0cdd0de50556737b14f..f4be333713af1e8e0428116373013536c2da8bab 100644
>--- a/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp
>+++ b/Source/WebCore/Modules/webgpu/WHLSL/Metal/WHLSLEntryPointScaffolding.cpp
>@@ -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 @@ String EntryPointScaffolding::resourceHelperTypes()
>     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 @@ String EntryPointScaffolding::resourceHelperTypes()
>             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();
>diff --git a/Source/WebCore/platform/graphics/gpu/cocoa/GPUBindGroupLayoutMetal.mm b/Source/WebCore/platform/graphics/gpu/cocoa/GPUBindGroupLayoutMetal.mm
>index 35ae3c96420185102c9368c4c60b55aedb1a4772..8020163554e8be4bb6115ad43f582c35fc5d4054 100644
>--- a/Source/WebCore/platform/graphics/gpu/cocoa/GPUBindGroupLayoutMetal.mm
>+++ b/Source/WebCore/platform/graphics/gpu/cocoa/GPUBindGroupLayoutMetal.mm
>@@ -96,7 +96,7 @@ RefPtr<GPUBindGroupLayout> GPUBindGroupLayout::tryCreate(const GPUDevice& device
>         return nullptr;
>     }
> 
>-    ArgumentArray vertexArgsArray, fragmentArgsArray, computeArgsArray;
>+    ArgumentArray vertexArgs, fragmentArgs, computeArgs, vertexLengths, fragmentLengths, computeLengths;
>     BindingsMapType bindingsMap;
> 
>     unsigned internalName = 0;
>@@ -137,33 +137,39 @@ RefPtr<GPUBindGroupLayout> GPUBindGroupLayout::tryCreate(const GPUDevice& device
>             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) {
>                     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)));
>diff --git a/LayoutTests/ChangeLog b/LayoutTests/ChangeLog
>index 28c62a6c22bfe01b183cd3f97c92a682306ef1d4..421ad08c645d7c43f8fba459a234f373d32c5d7e 100644
>--- a/LayoutTests/ChangeLog
>+++ b/LayoutTests/ChangeLog
>@@ -1,3 +1,47 @@
>+2019-06-18  Justin Fan  <justin_fan@apple.com>
>+
>+        [WHLSL] Create a shading language test harness
>+        https://bugs.webkit.org/show_bug.cgi?id=198978
>+
>+        Reviewed by NOBODY (OOPS!).
>+
>+        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-18  Russell Epstein  <russell_e@apple.com>
> 
>         Layout Test imported/w3c/web-platform-tests/content-security-policy/reporting/report-only-in-meta.sub.html is failing.
>diff --git a/LayoutTests/webgpu/js/whlsl-test-harness.js b/LayoutTests/webgpu/js/whlsl-test-harness.js
>new file mode 100644
>index 0000000000000000000000000000000000000000..2e85b88fc76e429d8f7eb02163eff8636d4210f6
>--- /dev/null
>+++ b/LayoutTests/webgpu/js/whlsl-test-harness.js
>@@ -0,0 +1,379 @@
>+/* 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 {
>+    _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 = 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 = 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
>diff --git a/LayoutTests/webgpu/whlsl-harness-test-expected.txt b/LayoutTests/webgpu/whlsl-harness-test-expected.txt
>new file mode 100644
>index 0000000000000000000000000000000000000000..e87095632c61eae31fb81dd47d0507d2b3d8bdc0
>--- /dev/null
>+++ b/LayoutTests/webgpu/whlsl-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/whlsl-harness-test.html b/LayoutTests/webgpu/whlsl-harness-test.html
>new file mode 100644
>index 0000000000000000000000000000000000000000..c00081d7834f7bc742fd37c20802aead0be79e85
>--- /dev/null
>+++ b/LayoutTests/webgpu/whlsl-harness-test.html
>@@ -0,0 +1,208 @@
>+<!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

Actions: View | Formatted Diff | Diff

Attachments on bug 198978: 372399 | 372407 | 372420 | 372425 | 372519