WebGPU: Compute shader encoders out-of-order

README.md

WebGPU: Compute shader encoders out-of-order

view on jsgist

index.css

pre { margin: 0; }

index.html

/*bug-in-github-api-content-can-not-be-empty*/

index.js

const arrayProd = arr => arr.reduce((a, b) => a * b);
const range = (num, fn) => new Array(num).fill(0).map((_, i) => fn(i));

const code = `
@group(0) @binding(0) var<storage, read_write> value: array<f32>;

@compute @workgroup_size(1) fn cs_add() { value[0] = value[0] + value[1]; }
@compute @workgroup_size(1) fn cs_sub() { value[0] = value[0] - value[1]; }
@compute @workgroup_size(1) fn cs_mul() { value[0] = value[0] * value[1]; }
@compute @workgroup_size(1) fn cs_div() { value[0] = value[0] / value[1]; }
`;

async function main(){
  const adapter = await navigator.gpu?.requestAdapter();
  const device = await adapter?.requestDevice();
  if (!device) {
    console.error('need WebGPU');
    return;
  }
  device.addEventListener('uncapturederror', e => console.error(e.error.message));

  const valueBuffer = device.createBuffer({ size: 8, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC | GPUBufferUsage.COPY_DST });
  device.queue.writeBuffer(valueBuffer, 0, new Float32Array([3, 5]));

  const shaderModule = device.createShaderModule({code});

  const bgl = device.createBindGroupLayout({
    entries: [
      {
        binding: 0,
        visibility: GPUShaderStage.COMPUTE,
        buffer: {
          type: "storage",
        },
      },
    ],
  });

  const layout = device.createPipelineLayout({
    bindGroupLayouts: [bgl],
  });

  const pipelines = ['add', 'sub', 'mul', 'div'].map(op => {
    return device.createComputePipeline({
      layout,
      compute: {
        module: shaderModule,
        entryPoint: `cs_${op}`,
      }
    });
  });

  const bindGroup = device.createBindGroup({
    layout: bgl,
    entries: [
      { binding: 0, resource: { buffer: valueBuffer }},
    ],
  });

  const add = 0;
  const sub = 1;
  const mul = 2;
  const div = 3;
  const cpy = 4;

  const cbBuffers = [
    [
      add, //  3 + 5 =  8, 5
      mul, //  8 * 5 = 40, 5
      sub, // 40 - 5 = 35, 5
      div, // 35 / 5 =  7, 5
      cpy,
    ],
    [
      sub, //  7, 5 =  2, 5
      mul, //  2, 5 = 10, 5
      add, // 10, 5 = 15, 5
      div, // 15, 5 =  3, 5
      cpy,
    ],
    [
      add, //   3, 5 =  8, 5
      div, //   8, 5 =  1.6, 5
      sub, // 1.6, 5 = -3.4, 5
      mul, //-3.4, 5 = -17, 5
      cpy,
    ],      
  ].map(ops => {
    const commandEncoders = pipelines.map(() => device.createCommandEncoder());
    const passEncoders = commandEncoders.map(enc => enc.beginComputePass());
    passEncoders.forEach((pass, i) => { pass.setPipeline(pipelines[i]); });
    passEncoders.forEach(pass => { pass.setBindGroup(0, bindGroup); });
    passEncoders.forEach(pass => { pass.dispatchWorkgroups(1); });
    passEncoders.forEach(pass => { pass.end(); });
    const commandBuffers = commandEncoders.map(enc => enc.finish());

    const resultBuffer = device.createBuffer({ size: valueBuffer.size, usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ });
    const commandEncoder = device.createCommandEncoder();
    commandEncoder.copyBufferToBuffer(valueBuffer, 0, resultBuffer, 0);
    commandBuffers.push(commandEncoder.finish());

    const cb = ops.map(i => commandBuffers[i]);

    return {cb, resultBuffer};
  });

  const cbs = cbBuffers.map(v => v.cb).flat();
  device.queue.submit(cbs);

  await Promise.all(cbBuffers.map(v => v.resultBuffer.mapAsync(GPUMapMode.READ)));

  cbBuffers.forEach(({resultBuffer}, i) => {
    const result = new Float32Array(resultBuffer.getMappedRange());
    console.log(`#${i}:`, result);
    resultBuffer.unmap();
  });
}

function log(...args) {
  const elem = document.createElement('pre');
  elem.textContent = args.join(' ');
  document.body.appendChild(elem);
}

main();

jsGist.json

{"name":"WebGPU: Compute shader encoders out-of-order","settings":{},"filenames":["index.html","index.css","index.js"]}