From 0944c8857067fe94145b6271797e3544577fe104 Mon Sep 17 00:00:00 2001 From: David Neto Date: Fri, 26 Jan 2024 13:52:48 -0500 Subject: [PATCH] wgsl: Add compound statement tests Validation: - Check parsing of compound statements - Update break-if statement to add a case for a plain compound statement surrounding a break-if. - For a declaration in a compound statement, check potential uses and potentially conflicting declarations. Execution: Add execution tests for compound statement. Observe a value at various points relative to its declaration inside a compound statement. Add timings for new tests Fixed: #1634 --- src/webgpu/listing_meta.json | 4 + .../execution/statement/compound.spec.ts | 137 ++++++++++++++++++ .../decl/compound_statement.spec.ts | 98 +++++++++++++ .../shader/validation/parse/break_if.spec.ts | 4 + .../shader/validation/parse/compound.spec.ts | 52 +++++++ 5 files changed, 295 insertions(+) create mode 100644 src/webgpu/shader/execution/statement/compound.spec.ts create mode 100644 src/webgpu/shader/validation/decl/compound_statement.spec.ts create mode 100644 src/webgpu/shader/validation/parse/compound.spec.ts diff --git a/src/webgpu/listing_meta.json b/src/webgpu/listing_meta.json index dcebf8fdadc7..2dc721e85989 100644 --- a/src/webgpu/listing_meta.json +++ b/src/webgpu/listing_meta.json @@ -1694,6 +1694,7 @@ "webgpu:shader,execution,shadow:while:*": { "subcaseMS": 7.400 }, "webgpu:shader,execution,stage:basic_compute:*": { "subcaseMS": 1.000 }, "webgpu:shader,execution,stage:basic_render:*": { "subcaseMS": 1.000 }, + "webgpu:shader,execution,statement,compound:decl:*": { "subcaseMS": 29.767 }, "webgpu:shader,execution,statement,increment_decrement:frexp_exp_increment:*": { "subcaseMS": 4.700 }, "webgpu:shader,execution,statement,increment_decrement:scalar_i32_decrement:*": { "subcaseMS": 20.301 }, "webgpu:shader,execution,statement,increment_decrement:scalar_i32_decrement_underflow:*": { "subcaseMS": 4.900 }, @@ -1717,6 +1718,8 @@ "webgpu:shader,validation,const_assert,const_assert:constant_expression_logical_or_no_assert:*": { "subcaseMS": 1.373 }, "webgpu:shader,validation,const_assert,const_assert:constant_expression_no_assert:*": { "subcaseMS": 1.655 }, "webgpu:shader,validation,const_assert,const_assert:evaluation_stage:*": { "subcaseMS": 3.367 }, + "webgpu:shader,validation,decl,compound_statement:decl_conflict:*": { "subcaseMS": 5.225 }, + "webgpu:shader,validation,decl,compound_statement:decl_use:*": { "subcaseMS": 0.625 }, "webgpu:shader,validation,decl,const:no_direct_recursion:*": { "subcaseMS": 0.951 }, "webgpu:shader,validation,decl,const:no_indirect_recursion:*": { "subcaseMS": 0.950 }, "webgpu:shader,validation,decl,const:no_indirect_recursion_via_array_size:*": { "subcaseMS": 2.601 }, @@ -1889,6 +1892,7 @@ "webgpu:shader,validation,parse,comments:line_comment_eof:*": { "subcaseMS": 4.500 }, "webgpu:shader,validation,parse,comments:line_comment_terminators:*": { "subcaseMS": 1.021 }, "webgpu:shader,validation,parse,comments:unterminated_block_comment:*": { "subcaseMS": 8.950 }, + "webgpu:shader,validation,parse,compound:parse:*": { "subcaseMS": 4.315 }, "webgpu:shader,validation,parse,const:placement:*": { "subcaseMS": 1.167 }, "webgpu:shader,validation,parse,const_assert:parse:*": { "subcaseMS": 1.400 }, "webgpu:shader,validation,parse,diagnostic:after_other_directives:*": { "subcaseMS": 1.000 }, diff --git a/src/webgpu/shader/execution/statement/compound.spec.ts b/src/webgpu/shader/execution/statement/compound.spec.ts new file mode 100644 index 000000000000..aed0cc224557 --- /dev/null +++ b/src/webgpu/shader/execution/statement/compound.spec.ts @@ -0,0 +1,137 @@ +export const description = ` +Compound statement execution. +`; + +import { makeTestGroup } from '../../../../common/framework/test_group.js'; +import { keysOf } from '../../../../common/util/data_tables.js'; +import { TypedArrayBufferView } from '../../../../common/util/util.js'; +import { GPUTest } from '../../../gpu_test.js'; + +export const g = makeTestGroup(GPUTest); + +/** + * Builds, runs then checks the output of a statement shader test. + * + * @param t The test object + * @param ty The WGSL scalar type to be written + * @param values The expected output values of type ty + * @param wgsl_main The body of the WGSL entry point. + */ +export function runStatementTest( + t: GPUTest, + ty: string, + values: TypedArrayBufferView, + wgsl_main: string +) { + const wgsl = ` +struct Outputs { + data : array<${ty}>, +}; +var count: u32 = 0; + +@group(0) @binding(1) var outputs : Outputs; + +fn put(value : ${ty}) { + outputs.data[count] = value; + count += 1; +} + +@compute @workgroup_size(1) +fn main() { + _ = &outputs; + ${wgsl_main} +} +`; + + const pipeline = t.device.createComputePipeline({ + layout: 'auto', + compute: { + module: t.device.createShaderModule({ code: wgsl }), + entryPoint: 'main', + }, + }); + + const maxOutputValues = 1000; + const outputBuffer = t.device.createBuffer({ + size: 4 * (1 + maxOutputValues), + usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC, + }); + + const bindGroup = t.device.createBindGroup({ + layout: pipeline.getBindGroupLayout(0), + entries: [{ binding: 1, resource: { buffer: outputBuffer } }], + }); + + // Run the shader. + const encoder = t.device.createCommandEncoder(); + const pass = encoder.beginComputePass(); + pass.setPipeline(pipeline); + pass.setBindGroup(0, bindGroup); + pass.dispatchWorkgroups(1); + pass.end(); + t.queue.submit([encoder.finish()]); + + t.expectGPUBufferValuesEqual(outputBuffer, values); +} + +// Consider a declaration X of identifier 'x' inside a compound statement. +// Check the value of 'x' at various places relative to X: +// a { b; X=c; d; { e; } } f; + +const kTests = { + uses: { + // Observe values without conflicting declarations. + src: `let x = 1; + put(x); + { + put(x); + let x = x+1; // The declaration in question + put(x); + { + put(x); + } + put(x); + } + put(x);`, + values: [1, 1, 2, 2, 2, 1], + }, + shadowed: { + // Observe values when shadowed + src: `let x = 1; + put(x); + { + put(x); + let x = x+1; // The declaration in question + put(x); + { + let x = x+1; // A shadow + put(x); + } + put(x); + } + put(x);`, + values: [1, 1, 2, 3, 2, 1], + }, + gone: { + // The declaration goes out of scope. + src: `{ + let x = 2; // The declaration in question + put(x); + } + let x = 1; + put(x);`, + values: [2, 1], + }, +} as const; + +g.test('decl') + .desc('Tests the value of a declared value in a compound statment.') + .params(u => u.combine('case', keysOf(kTests))) + .fn(t => { + runStatementTest( + t, + 'i32', + new Int32Array(kTests[t.params.case].values), + kTests[t.params.case].src + ); + }); diff --git a/src/webgpu/shader/validation/decl/compound_statement.spec.ts b/src/webgpu/shader/validation/decl/compound_statement.spec.ts new file mode 100644 index 000000000000..8ad89f48a8aa --- /dev/null +++ b/src/webgpu/shader/validation/decl/compound_statement.spec.ts @@ -0,0 +1,98 @@ +export const description = ` +Validation tests for declarations in compound statements. +`; + +import { makeTestGroup } from '../../../../common/framework/test_group.js'; +import { keysOf } from '../../../../common/util/data_tables.js'; +import { ShaderValidationTest } from '../shader_validation_test.js'; + +export const g = makeTestGroup(ShaderValidationTest); + +// 9.1 Compound Statements +// When a declaration is one of those statements, its identifier is in scope from +// the start of the next statement until the end of the compound statement. +// +// Enumerating cases: Consider a declaration X inside a compound statement. +// The X declaration should be tested with potential uses and potentially +// conflicting declarations in positions [a, b, c, d, e], in the following: +// a { b; X; c; { d; } } e; + +const kConflictTests = { + a: { + src: 'let x = 1; { let x = 1; }', + pass: true, + }, + bc: { + src: '{let x = 1; let x = 1; }', + pass: false, + }, + d: { + src: '{let x = 1; { let x = 1; }}', + pass: true, + }, + e: { + src: '{let x = 1; } let x = 1;', + pass: true, + }, +}; + +g.test('decl_conflict') + .desc( + 'Test a potentially conflicting declaration relative to a declaration in a compound statement' + ) + .params(u => u.combine('case', keysOf(kConflictTests))) + .fn(t => { + const wgsl = ` +@vertex fn vtx() -> @builtin(position) vec4f { + ${kConflictTests[t.params.case].src} + return vec4f(1); +}`; + t.expectCompileResult(kConflictTests[t.params.case].pass, wgsl); + }); + +const kUseTests = { + a: { + src: 'let y = x; { let x = 1; }', + pass: false, // not visible + }, + b: { + src: '{ let y = x; let x = 1; }', + pass: false, // not visible + }, + self: { + src: '{ let x = (x);}', + pass: false, // not visible + }, + c_yes: { + src: '{ const x = 1; const_assert x == 1; }', + pass: true, + }, + c_no: { + src: '{ const x = 1; const_assert x == 2; }', + pass: false, + }, + d_yes: { + src: '{ const x = 1; { const_assert x == 1; }}', + pass: true, + }, + d_no: { + src: '{ const x = 1; { const_assert x == 2; }}', + pass: false, + }, + e: { + src: '{ const x = 1; } let y = x;', + pass: false, // not visible + }, +}; + +g.test('decl_use') + .desc('Test a use of a declaration in a compound statement') + .params(u => u.combine('case', keysOf(kUseTests))) + .fn(t => { + const wgsl = ` +@vertex fn vtx() -> @builtin(position) vec4f { + ${kUseTests[t.params.case].src} + return vec4f(1); +}`; + t.expectCompileResult(kUseTests[t.params.case].pass, wgsl); + }); diff --git a/src/webgpu/shader/validation/parse/break_if.spec.ts b/src/webgpu/shader/validation/parse/break_if.spec.ts index 41d448437e45..97a625f625cc 100644 --- a/src/webgpu/shader/validation/parse/break_if.spec.ts +++ b/src/webgpu/shader/validation/parse/break_if.spec.ts @@ -7,6 +7,10 @@ import { ShaderValidationTest } from '../shader_validation_test.js'; export const g = makeTestGroup(ShaderValidationTest); const kTests = { + compound_break: { + src: '{ break if true; }', + pass: false, + }, loop_break: { src: 'loop { break if true; }', pass: false, diff --git a/src/webgpu/shader/validation/parse/compound.spec.ts b/src/webgpu/shader/validation/parse/compound.spec.ts new file mode 100644 index 000000000000..b3627c2e5b9e --- /dev/null +++ b/src/webgpu/shader/validation/parse/compound.spec.ts @@ -0,0 +1,52 @@ +export const description = `Validation tests for compound statements`; + +import { makeTestGroup } from '../../../../common/framework/test_group.js'; +import { keysOf } from '../../../../common/util/data_tables.js'; +import { ShaderValidationTest } from '../shader_validation_test.js'; + +export const g = makeTestGroup(ShaderValidationTest); + +const kTests = { + missing_start: { + src: '}', + pass: false, + }, + missing_end: { + src: '{', + pass: false, + }, + empty: { + src: '{}', + pass: true, + }, + semicolon: { + src: '{;}', + pass: true, + }, + semicolons: { + src: '{;;}', + pass: true, + }, + decl: { + src: '{const c = 1;}', + pass: true, + }, + nested: { + src: '{ {} }', + pass: true, + }, +}; + +g.test('parse') + .desc('Test that compound statments parse') + .params(u => u.combine('stmt', keysOf(kTests))) + .fn(t => { + const code = ` +@vertex +fn vtx() -> @builtin(position) vec4f { + ${kTests[t.params.stmt].src} + return vec4f(1); +} + `; + t.expectCompileResult(kTests[t.params.stmt].pass, code); + });