import { Colors } from '../../../../common/util/colors.js'; import { AllFeaturesMaxLimitsGPUTest } from '../../../gpu_test.js'; /** * Options for runFlowControlTest() */ interface FlowControlTest extends AllFeaturesMaxLimitsGPUTest { params: { /** * If true, then constant values will be placed into a storage buffer, * preventing the shader compiler from knowing the value at compile time. * This can prevent constant folding, loop unrolling, dead-code * optimizations etc, which would could all affect the tests. */ preventValueOptimizations?: boolean; }; } /** * The builder interface for the runFlowControlTest() callback. * This interface is indented to be used to inject WGSL logic into the test * shader. * @see runFlowControlTest */ interface FlowControlTestBuilder { /** * Emits a value into the shader. * If the FlowControlTest.params.preventValueOptimizations flag is enabled, * then value() emits an expression to load the given value from a storage * buffer, preventing the shader compiler from knowing the value at compile * time. This can prevent constant folding, loop unrolling, dead-code * optimizations etc, which would could all affect the tests. */ value(v: number | boolean): string; /** * Emits an expectation that the statement will be executed at the given * chronological events. * @param event one or more chronological events, the first being 0. */ expect_order(...event: number[]): string; /** * Emits an expectation that the statement will not be reached. */ expect_not_reached(): string; } /** * Builds, runs then checks the output of a flow control shader test. * * `build_wgsl` is a function that's called to build the WGSL shader. * This function takes a FlowControlTestBuilder as the single argument, and * returns either a string which is embedded into the WGSL entrypoint function, * or an object of the signature `{ entrypoint: string; extra: string }` which * contains the entrypoint code, along with additional module-scope code. * * The FlowControlTestBuilder should be used to insert expectations into WGSL to * validate control flow. FlowControlTestBuilder also can be used to add values * to the shader which cannot be optimized away. * * Example, testing that an if-statement behaves as expected: * * ``` * runFlowControlTest(t, f => * ` * ${f.expect_order(0)} * if (${f.value(true)}) { * ${f.expect_order(1)} * } else { * ${f.expect_not_reached()} * } * ${f.expect_order(2)} * `); * ``` * * @param t The test object * @param builder The shader builder function that takes a * FlowControlTestBuilder as the single argument, and returns either a WGSL * string which is embedded into the WGSL entrypoint function, or a structure * with entrypoint-scoped WGSL code and extra module-scope WGSL code. */ export function runFlowControlTest( t: FlowControlTest, build_wgsl: (builder: FlowControlTestBuilder) => string | { entrypoint: string; extra: string } ) { const inputData = new Array(); type ExpectedEvents = { kind: 'events'; stack: string | undefined; values: number[]; counter: number; }; type ExpectedNotReached = { kind: 'not-reached'; stack: string | undefined; }; const expectations = new Array(); const build_wgsl_result = build_wgsl({ value: v => { if (t.params.preventValueOptimizations) { if (typeof v === 'boolean') { inputData.push(v ? 1 : 0); return `inputs[${inputData.length - 1}] != 0`; } inputData.push(v); return `inputs[${inputData.length - 1}]`; } else { return `${v}`; } }, expect_order: (...expected) => { expectations.push({ kind: 'events', stack: Error().stack, values: expected, counter: 0, }); // Expectation id starts from 1 to distinguish from initialization 0. return `push_output(${expectations.length}); // expect_order(${expected.join(', ')})`; }, expect_not_reached: () => { expectations.push({ kind: 'not-reached', stack: Error().stack, }); // Expectation id starts from 1 to distinguish from initialization 0. return `push_output(${expectations.length}); // expect_not_reached()`; }, }); const built_wgsl = typeof build_wgsl_result === 'string' ? { entrypoint: build_wgsl_result, extra: '' } : build_wgsl_result; const main_wgsl = built_wgsl.entrypoint !== undefined ? built_wgsl : built_wgsl.entrypoint; const wgsl = ` struct Outputs { count : u32, data : array, }; @group(0) @binding(0) var inputs : array; @group(0) @binding(1) var outputs : Outputs; fn push_output(value : u32) { outputs.data[outputs.count] = value; outputs.count++; } @compute @workgroup_size(1) fn main() { _ = &inputs; _ = &outputs; ${main_wgsl.entrypoint} } ${main_wgsl.extra} `; const pipeline = t.device.createComputePipeline({ layout: 'auto', compute: { module: t.device.createShaderModule({ code: wgsl }), entryPoint: 'main', }, }); // If there are no inputs, just put a single value in the buffer to keep // makeBufferWithContents() happy. if (inputData.length === 0) { inputData.push(0); } const inputBuffer = t.makeBufferWithContents(new Uint32Array(inputData), GPUBufferUsage.STORAGE); const maxOutputValues = 1000; const outputBuffer = t.createBufferTracked({ size: 4 * (1 + maxOutputValues), usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC, }); const bindGroup = t.device.createBindGroup({ layout: pipeline.getBindGroupLayout(0), entries: [ { binding: 0, resource: { buffer: inputBuffer } }, { binding: 1, resource: { buffer: outputBuffer } }, ], }); // Run the shader. const encoder = t.device.createCommandEncoder({ label: 'runFlowControlTest' }); const pass = encoder.beginComputePass(); pass.setPipeline(pipeline); pass.setBindGroup(0, bindGroup); pass.dispatchWorkgroups(1); pass.end(); t.queue.submit([encoder.finish()]); t.eventualExpectOK( t .readGPUBufferRangeTyped(outputBuffer, { type: Uint32Array, typedLength: outputBuffer.size / 4, }) .then(outputs => { // outputs[0] is the number of outputted values // outputs[1..N] holds the outputted values const outputCount = outputs.data[0]; if (outputCount > maxOutputValues) { return new Error( `output data count (${outputCount}) exceeds limit of ${maxOutputValues}` ); } // returns an Error with the given message and WGSL source const fail = (err: string) => Error(`${err}\nWGSL:\n${Colors.dim(Colors.blue(wgsl))}`); // returns a string that shows the outputted values to help understand the whole trace. const print_output_value = () => { const subarray = outputs.data.subarray(1, outputCount + 1); return `Output values (length: ${outputCount}): ${subarray.join(', ')}`; }; // returns a colorized string of the expect_order() call, highlighting // the event number that caused an error. const expect_order_err = (expectation: ExpectedEvents, err_idx: number) => { let out = 'expect_order('; for (let i = 0; i < expectation.values.length; i++) { if (i > 0) { out += ', '; } if (i < err_idx) { out += Colors.green(`${expectation.values[i]}`); } else if (i > err_idx) { out += Colors.dim(`${expectation.values[i]}`); } else { out += Colors.red(`${expectation.values[i]}`); } } out += ')'; return out; }; // Each of the outputted values represents an event // Check that each event is as expected for (let event = 0; event < outputCount; event++) { const eventValue = outputs.data[1 + event]; // outputs.data[0] is count // Expectation id starts from 1, and 0 is invalid value. if (eventValue === 0) { return fail( `outputs.data[${event}] is initial value 0, doesn't refer to any valid expectations)\n${print_output_value()}` ); } const expectationIndex = eventValue - 1; if (expectationIndex >= expectations.length) { return fail( `outputs.data[${event}] value (${expectationIndex}) exceeds number of expectations (${ expectations.length })\n${print_output_value()}` ); } const expectation = expectations[expectationIndex]; switch (expectation.kind) { case 'not-reached': return fail( `expect_not_reached() reached at event ${event}\n${print_output_value()}\n${ expectation.stack }` ); case 'events': if (expectation.counter >= expectation.values.length) { return fail( `${expect_order_err( expectation, expectation.counter )}) unexpectedly reached at event ${Colors.red( `${event}` )}\n${print_output_value()}\n${expectation.stack}` ); } if (event !== expectation.values[expectation.counter]) { return fail( `${expect_order_err(expectation, expectation.counter)} expected event ${ expectation.values[expectation.counter] }, got ${event}\n${print_output_value()}\n${expectation.stack}` ); } expectation.counter++; break; } } // Finally check that all expect_order() calls were reached for (const expectation of expectations) { if (expectation.kind === 'events' && expectation.counter !== expectation.values.length) { return fail( `${expect_order_err(expectation, expectation.counter)} event ${ expectation.values[expectation.counter] } was not reached\n${expectation.stack}\n${print_output_value()}` ); } } outputs.cleanup(); return undefined; }) ); }