wgsl: Add compound statement tests (#3337)

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

@@ -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 },

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<private> count: u32 = 0;
+
+@group(0) @binding(1) var<storage, read_write> 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
+    );
+  });

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);
+  });

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,

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);
+  });