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

export const description = `
Increment and decrement statement tests.
`;

import { makeTestGroup } from '../../../../common/framework/test_group.js';
import { TypedArrayBufferView } from '../../../../common/util/util.js';
import { AllFeaturesMaxLimitsGPUTest, GPUTest } from '../../../gpu_test.js';
import { kValue } from '../../../util/constants.js';

export const g = makeTestGroup(AllFeaturesMaxLimitsGPUTest);

/**
 * Builds, runs then checks the output of a statement shader test.
 *
 * @param t The test object
 * @param builder The shader builder function that takes a
 * StatementTestBuilder 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 runStatementTest(
  t: GPUTest,
  fmt: string,
  values: TypedArrayBufferView,
  wgsl_main: string,
  extras: { global_decl?: string } = {}
) {
  const wgsl = `
struct Outputs {
  data  : array<${fmt}>,
};
var<private> count: u32 = 0;

@group(0) @binding(1) var<storage, read_write> outputs : Outputs;

fn push_output(value : ${fmt}) {
  outputs.data[count] = value;
  count += 1;
}

${extras.global_decl ?? ''}

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

g.test('scalar_i32_increment')
  .desc('Tests increment of scalar i32 values')
  .fn(t => {
    runStatementTest(
      t,
      'i32',
      new Int32Array([-9, 11, kValue.i32.negative.min + 1, kValue.i32.positive.max, 1]),
      `
    var a: i32 = -10;
    var b: i32 = 10;
    var c: i32 = ${kValue.i32.negative.min};
    var d: i32 = ${kValue.i32.positive.max - 1};
    var e: i32 = 0;

    a++;
    b++;
    c++;
    d++;
    e++;

    push_output(a);
    push_output(b);
    push_output(c);
    push_output(d);
    push_output(e);
`
    );
  });

g.test('scalar_i32_increment_overflow')
  .desc('Tests increment of scalar i32 values which overflows')
  .fn(t => {
    runStatementTest(
      t,
      'i32',
      new Int32Array([kValue.i32.negative.min]),
      `
    var a: i32 = ${kValue.i32.positive.max};
    a++;
    push_output(a);
`
    );
  });

g.test('scalar_u32_increment')
  .desc('Tests increment of scalar u32 values')
  .fn(t => {
    runStatementTest(
      t,
      'u32',
      new Uint32Array([1, 11, kValue.u32.max]),
      `
    var a: u32 = 0;
    var b: u32 = 10;
    var c: u32 = ${kValue.u32.max - 1};

    a++;
    b++;
    c++;

    push_output(a);
    push_output(b);
    push_output(c);
`
    );
  });

g.test('scalar_u32_increment_overflow')
  .desc('Tests increment of scalar u32 values which overflows')
  .fn(t => {
    runStatementTest(
      t,
      'u32',
      new Uint32Array([0]),
      `
    var a: u32 = ${kValue.u32.max};
    a++;
    push_output(a);
`
    );
  });

g.test('scalar_i32_decrement')
  .desc('Tests decrement of scalar i32 values')
  .fn(t => {
    runStatementTest(
      t,
      'i32',
      new Int32Array([-11, 9, kValue.i32.negative.min, kValue.i32.positive.max - 1, -1]),
      `
    var a: i32 = -10;
    var b: i32 = 10;
    var c: i32 = ${kValue.i32.negative.min + 1};
    var d: i32 = ${kValue.i32.positive.max};
    var e: i32 = 0;

    a--;
    b--;
    c--;
    d--;
    e--;

    push_output(a);
    push_output(b);
    push_output(c);
    push_output(d);
    push_output(e);
`
    );
  });

g.test('scalar_i32_decrement_underflow')
  .desc('Tests decrement of scalar i32 values which underflow')
  .fn(t => {
    runStatementTest(
      t,
      'i32',
      new Int32Array([kValue.i32.positive.max]),
      `
    var a: i32 = ${kValue.i32.negative.min};
    a--;
    push_output(a);
`
    );
  });

g.test('scalar_u32_decrement')
  .desc('Tests decrement of scalar u32 values')
  .fn(t => {
    runStatementTest(
      t,
      'u32',
      new Uint32Array([0, 9, kValue.u32.max - 1]),
      `
    var a: u32 = 1;
    var b: u32 = 10;
    var c: u32 = ${kValue.u32.max};

    a--;
    b--;
    c--;

    push_output(a);
    push_output(b);
    push_output(c);
`
    );
  });

g.test('scalar_u32_decrement_underflow')
  .desc('Tests decrement of scalar u32 values which underflow')
  .fn(t => {
    runStatementTest(
      t,
      'u32',
      new Uint32Array([kValue.u32.max]),
      `
    var a: u32 = 0;
    a--;
    push_output(a);
`
    );
  });

g.test('vec2_element_increment')
  .desc('Tests increment of ve2 values')
  .fn(t => {
    runStatementTest(
      t,
      'i32',
      new Int32Array([-9, 11]),
      `
    var a = vec2(-10, 10);

    a.x++;
    a.g++;

    push_output(a.x);
    push_output(a.y);
`
    );
  });

g.test('vec3_element_increment')
  .desc('Tests increment of vec3 values')
  .fn(t => {
    runStatementTest(
      t,
      'i32',
      new Int32Array([-9, 11, kValue.i32.negative.min + 1]),
      `
    var a = vec3(-10, 10, ${kValue.i32.negative.min});

    a.x++;
    a.g++;
    a.z++;

    push_output(a.x);
    push_output(a.y);
    push_output(a.z);
`
    );
  });

g.test('vec4_element_increment')
  .desc('Tests increment of vec4 values')
  .fn(t => {
    runStatementTest(
      t,
      'i32',
      new Int32Array([-9, 11, kValue.i32.negative.min + 1, kValue.i32.positive.max]),
      `
    var a: vec4<i32> = vec4(-10, 10, ${kValue.i32.negative.min}, ${kValue.i32.positive.max - 1});

    a.x++;
    a.g++;
    a.z++;
    a.a++;

    push_output(a.x);
    push_output(a.y);
    push_output(a.z);
    push_output(a.w);
`
    );
  });

g.test('vec2_element_decrement')
  .desc('Tests decrement of vec2 values')
  .fn(t => {
    runStatementTest(
      t,
      'i32',
      new Int32Array([-11, 9]),
      `
    var a = vec2(-10, 10);

    a.x--;
    a.g--;

    push_output(a.x);
    push_output(a.y);
`
    );
  });

g.test('vec3_element_decrement')
  .desc('Tests decrement of vec3 values')
  .fn(t => {
    runStatementTest(
      t,
      'i32',
      new Int32Array([-11, 9, kValue.i32.negative.min]),
      `
    var a = vec3(-10, 10, ${kValue.i32.negative.min + 1});

    a.x--;
    a.g--;
    a.z--;

    push_output(a.x);
    push_output(a.y);
    push_output(a.z);
`
    );
  });

g.test('vec4_element_decrement')
  .desc('Tests decrement of vec4 values')
  .fn(t => {
    runStatementTest(
      t,
      'i32',
      new Int32Array([-11, 9, kValue.i32.negative.min, kValue.i32.positive.max - 1]),
      `
    var a: vec4<i32> = vec4(-10, 10, ${kValue.i32.negative.min + 1}, ${kValue.i32.positive.max});

    a.x--;
    a.g--;
    a.z--;
    a.a--;

    push_output(a.x);
    push_output(a.y);
    push_output(a.z);
    push_output(a.w);
`
    );
  });

g.test('frexp_exp_increment')
  .desc('Tests increment can be used on a frexp field')
  .fn(t => {
    runStatementTest(
      t,
      'i32',
      new Int32Array([2]),
      `
    var a = frexp(1.23);
    a.exp++;
    push_output(a.exp);
`
    );
  });

g.test('single_eval_increment')
  .desc('Tests the left-hand-side reference of an increment is computed only once.')
  .fn(t => {
    runStatementTest(
      t,
      'i32',
      new Int32Array([999, 0, 1, 2, 11, 21, 31]),
      `
    var a: array<i32,3> = array(10, 20, 30);

    push_output(999);
    a[bump()]++;
    a[bump()]++;
    a[bump()]++;
    push_output(a[0]);
    push_output(a[1]);
    push_output(a[2]);
`,
      {
        global_decl: `
  var<private> index_counter: i32 = 0;
  fn bump() -> i32 {
    let result = index_counter;
    push_output(result);
    index_counter = index_counter + 1;
    return result;
  }
  `,
      }
    );
  });

g.test('single_eval_decrement')
  .desc('Tests the left-hand-side reference of a decrement is computed only once.')
  .fn(t => {
    runStatementTest(
      t,
      'i32',
      new Int32Array([999, 0, 1, 2, 9, 19, 29]),
      `
    var a: array<i32,3> = array(10, 20, 30);

    push_output(999);
    a[bump()]--;
    a[bump()]--;
    a[bump()]--;
    push_output(a[0]);
    push_output(a[1]);
    push_output(a[2]);
`,
      {
        global_decl: `
  var<private> index_counter: i32 = 0;
  fn bump() -> i32 {
    let result = index_counter;
    push_output(result);
    index_counter = index_counter + 1;
    return result;
  }
  `,
      }
    );
  });