WebKit Bugzilla
Attachment 372420 Details for
Bug 198978
: [WHLSL] Create a shading language test harness
Home
|
New
|
Browse
|
Search
|
[?]
|
Reports
|
Requests
|
Help
|
New Account
|
Log In
Remember
[x]
|
Forgot Password
Login:
[x]
[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:
Patch
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><!-- webkit-test-runner [ experimental:WebGPUEnabled=true ] --> >+<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
You cannot view the attachment while viewing its details because your browser does not support IFRAMEs.
View the attachment on a separate page
.
View Attachment As Diff
View Attachment As Raw
Actions:
View
|
Formatted Diff
|
Diff
Attachments on
bug 198978
:
372399
|
372407
|
372420
|
372425
|
372519