<script src="../../resources/js-test-pre.js"></script>
<script>
  globalThis.testRunner?.waitUntilDone();
  const log = globalThis.$vm?.print ?? console.log;

  onload = async () => {
    let adapter = await navigator.gpu.requestAdapter({});
    let device = await adapter.requestDevice({});
    device.pushErrorScope('validation');
    let code = `
@group(0) @binding(0) var<storage, read_write> buf0: array<vec4u, 16>;

var<workgroup> vb0: bool;
var<workgroup> vb1: bool;
var<workgroup> vb2: bool;
var<workgroup> vb3: bool;
var<workgroup> vu: vec4u;

@compute @workgroup_size(1)
fn c() {
  var d0 = u32(vb0);
  var d1 = u32(vb1);
  var d2 = u32(vb2);
  var d3 = u32(vb3);
  if d0+d1+d2+d3 < 5 {
    buf0[2] = vec4(d0+d1+d2+d3);
  }
  buf0[d0+d1+d2+d3] = vu;
  buf0[0].x = d0;
  buf0[0].y = d1;
  buf0[0].z = d2;
  buf0[0].w = d3;

  vb0 = bool(buf0[0][0]);
  vb1 = bool(buf0[0][0]);
  vb2 = bool(buf0[0][0]);
  vb3 = bool(buf0[0][0]);
  vu = buf0[0];
}
`;
    let module = device.createShaderModule({code});
    let bindGroupLayout0 = device.createBindGroupLayout({
      entries: [
        {binding: 0, buffer: {type: 'storage'}, visibility: GPUShaderStage.COMPUTE},
      ],
    });
    let buffer0 = device.createBuffer({
      size: 256, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC,
    });
    await device.queue.onSubmittedWorkDone();
    let bindGroup0 = device.createBindGroup({
      layout: bindGroupLayout0, entries: [
        {binding: 0, resource: {buffer: buffer0}},
      ],
    });
    let pipelineLayout = device.createPipelineLayout({bindGroupLayouts: [bindGroupLayout0]});
    let commandEncoder = device.createCommandEncoder();
    let computePassEncoder = commandEncoder.beginComputePass({});
    let computePipeline = device.createComputePipeline({layout: pipelineLayout, compute: {module}});
    computePassEncoder.setPipeline(computePipeline);
    computePassEncoder.setBindGroup(0, bindGroup0);
    computePassEncoder.dispatchWorkgroups(1);
    computePassEncoder.end();
    let outputBuffer0 = device.createBuffer({size: buffer0.size, usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ});
    commandEncoder.copyBufferToBuffer(buffer0, 0, outputBuffer0, 0, buffer0.size);
    device.queue.submit([commandEncoder.finish()]);
    await device.queue.onSubmittedWorkDone();
    await outputBuffer0.mapAsync(GPUMapMode.READ);
    for (x of [...new Uint32Array(outputBuffer0.getMappedRange())]) {
      shouldBe('x', '0');
    }
    outputBuffer0.unmap();
    let error = await device.popErrorScope();
    if (error) {
      log(error.message);
    } else {
      log('no validation error');
    }
    globalThis.testRunner?.notifyDone();
  };
</script>