🏡 index : cold-air-networks/mirror-webgpu-cts.git

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 { AllFeaturesMaxLimitsGPUTest, GPUTest } from '../../../gpu_test.js';
import { runFlowControlTest } from '../flow_control/harness.js';

export const g = makeTestGroup(AllFeaturesMaxLimitsGPUTest);

/**
 * 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.createBufferTracked({
    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
    );
  });

g.test('eval_order')
  .desc('Tests evaluation order of compound assignment, lhs is evaluated before rhs')
  .fn(t => {
    runFlowControlTest(t, f => ({
      entrypoint: `
arr[0] = 41;
${f.expect_order(0)}
arr[idx()] += foo();
${f.expect_order(3)}
if (arr[0] == 42) {
  ${f.expect_order(4)}
} else {
  ${f.expect_not_reached()}
}
`,
      extra: `
var<private> arr : array<u32, 1>;
fn idx() -> u32 {
  ${f.expect_order(1)}
  return 0;
}
fn foo() -> u32 {
  ${f.expect_order(2)}
  arr[0] = 10;
  return 1;
}
`,
    }));
  });