export function runStorageVariableTest()

in src/webgpu/shader/execution/expression/call/builtin/atomics/harness.ts [39:112]

• 56 lines of code
• 3 McCabe index (conditional complexity)

export function runStorageVariableTest({
  t,
  workgroupSize, // Workgroup X-size
  dispatchSize, // Dispatch X-size
  bufferNumElements, // Number of 32-bit elements in output buffer
  initValue, // 32-bit initial value used to fill output buffer
  // Atomic op source executed by the compute shader, NOTE: 'id' is global_invocation_id.x,
  // and `output` is a storage array of atomics.
  op,
  expected, // Expected values array to compare against output buffer
  extra, // Optional extra WGSL source
}: {
  t: GPUTest;
  workgroupSize: number;
  dispatchSize: number;
  bufferNumElements: number;
  initValue: number;
  op: string;
  expected: TypedArrayBufferView;
  extra?: string;
}) {
  assert(expected.length === bufferNumElements, "'expected' buffer size is incorrect");

  const scalarType = expected instanceof Uint32Array ? 'u32' : 'i32';
  const arrayType = typedArrayCtor(scalarType);

  const wgsl = `
    @group(0) @binding(0)
    var<storage, read_write> output : array<atomic<${scalarType}>>;

    @compute @workgroup_size(${workgroupSize})
    fn main(
        @builtin(global_invocation_id) global_invocation_id : vec3<u32>,
        ) {
      let id = ${scalarType}(global_invocation_id[0]);
      ${op};
    }
    ${extra || ''}
    `;

  const pipeline = t.device.createComputePipeline({
    layout: 'auto',
    compute: {
      module: t.device.createShaderModule({ code: wgsl }),
      entryPoint: 'main',
    },
  });

  const outputBuffer = t.createBufferTracked({
    size: bufferNumElements * arrayType.BYTES_PER_ELEMENT,
    usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC,
    mappedAtCreation: true,
  });
  // Fill with initial value
  const data = new arrayType(outputBuffer.getMappedRange());
  data.fill(initValue);
  outputBuffer.unmap();

  const bindGroup = t.device.createBindGroup({
    layout: pipeline.getBindGroupLayout(0),
    entries: [{ binding: 0, resource: { buffer: outputBuffer } }],
  });

  // Run the shader.
  const encoder = t.device.createCommandEncoder({ label: 'runStorageVariableTest' });
  const pass = encoder.beginComputePass();
  pass.setPipeline(pipeline);
  pass.setBindGroup(0, bindGroup);
  pass.dispatchWorkgroups(dispatchSize);
  pass.end();
  t.queue.submit([encoder.finish()]);

  t.expectGPUBufferValuesEqual(outputBuffer, expected);
}