WebGPU Multiple Timestamps vs 1

README.md

WebGPU Multiple Timestamps vs 1

view on jsgist

index.css

@import url(https://webgpufundamentals.org/webgpu/resources/webgpu-lesson.css);

index.html

<div id="warn" style="display: none">feature: 'timestamp-query' is not available.<br>If you're using chrome pass in --enable-dawn-features=allow_unsafe_apis</div>
<pre id="info"></pre>
<div id="fail" style="display: none">
  <div class="content"></div>
</div>

index.js

function assert(cond, msg = '') {
  if (!cond) {
    throw new Error(msg);
  }
}

// We can't use a WeakSet because the command buffer might have been GCed
// between the time we submit and the time we call getResult.
const s_unsubmittedCommandBuffer = new Set();

/* global GPUQueue */
GPUQueue.prototype.submit = (function(origFn) {
  return function(commandBuffers) {
    origFn.call(this, commandBuffers);
    commandBuffers.forEach(cb => s_unsubmittedCommandBuffer.delete(cb));
  };
})(GPUQueue.prototype.submit);

// See https://webgpufundamentals.org/webgpu/lessons/webgpu-timing.html
export default class TimingHelper {
  #canTimestamp;
  #device;
  #querySet;
  #resolveBuffer;
  #resultBuffer;
  #commandBuffer;
  #resultBuffers = [];
  // state can be 'free', 'need resolve', 'wait for result'
  #state = 'free';

  constructor(device) {
    this.#device = device;
    this.#canTimestamp = device.features.has('timestamp-query');
    if (this.#canTimestamp) {
      this.#querySet = device.createQuerySet({
         type: 'timestamp',
         count: 2,
      });
      this.#resolveBuffer = device.createBuffer({
        size: this.#querySet.count * 8,
        usage: GPUBufferUsage.QUERY_RESOLVE | GPUBufferUsage.COPY_SRC,
      });
    }
  }

  #beginTimestampPass(encoder, fnName, descriptor) {
    if (this.#canTimestamp) {
      assert(this.#state === 'free', 'state not free');
      this.#state = 'need resolve';

      const pass = encoder[fnName]({
        ...descriptor,
        ...{
          timestampWrites: {
            querySet: this.#querySet,
            beginningOfPassWriteIndex: 0,
            endOfPassWriteIndex: 1,
          },
        },
      });

      const resolve = () => this.#resolveTiming(encoder);
      const trackCommandBuffer = (cb) => this.#trackCommandBuffer(cb);
      pass.end = (function(origFn) {
        return function() {
          origFn.call(this);
          resolve();
        };
      })(pass.end);

      encoder.finish = (function(origFn) {
        return function() {
          const cb = origFn.call(this);
          trackCommandBuffer(cb);
          return cb;
        };
      })(encoder.finish);

      return pass;
    } else {
      return encoder[fnName](descriptor);
    }
  }

  beginRenderPass(encoder, descriptor = {}) {
    return this.#beginTimestampPass(encoder, 'beginRenderPass', descriptor);
  }

  beginComputePass(encoder, descriptor = {}) {
    return this.#beginTimestampPass(encoder, 'beginComputePass', descriptor);
  }

  #trackCommandBuffer(cb) {
    if (!this.#canTimestamp) {
      return;
    }
    assert(this.#state === 'need finish', 'must call encoder.finish');
    this.#commandBuffer = cb;
    s_unsubmittedCommandBuffer.add(cb);
    this.#state = 'wait for result';
  }

  #resolveTiming(encoder) {
    if (!this.#canTimestamp) {
      return;
    }
    assert(this.#state === 'need resolve', 'must call addTimestampToPass');
    this.#state = 'need finish';

    this.#resultBuffer = this.#resultBuffers.pop() || this.#device.createBuffer({
      size: this.#resolveBuffer.size,
      usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ,
    });

    encoder.resolveQuerySet(this.#querySet, 0, this.#querySet.count, this.#resolveBuffer, 0);
    encoder.copyBufferToBuffer(this.#resolveBuffer, 0, this.#resultBuffer, 0, this.#resultBuffer.size);
  }

  async getResult() {
    if (!this.#canTimestamp) {
      return 0;
    }
    assert(this.#state === 'wait for result', 'must call resolveTiming');
    assert(!!this.#commandBuffer); // internal check
    assert(!s_unsubmittedCommandBuffer.has(this.#commandBuffer), 'you must submit the command buffer before you can read the result');
    this.#commandBuffer = undefined;
    this.#state = 'free';

    const resultBuffer = this.#resultBuffer;
    await resultBuffer.mapAsync(GPUMapMode.READ);
    const times = new BigInt64Array(resultBuffer.getMappedRange());
    const duration = Number(times[1] - times[0]);
    resultBuffer.unmap();
    this.#resultBuffers.push(resultBuffer);
    return duration;
  }
}

async function main() {
  const adapter = await navigator.gpu?.requestAdapter();
  const canTimestamp = adapter.features.has('timestamp-query');
  const device = await adapter?.requestDevice({requiredFeatures: [
    'timestamp-query',
  ]});
  if (!device) {
    fail('need a browser that supports WebGPU');
    return;
  }

  if (!canTimestamp) {
    document.querySelector('#warn').style.display = '';
  }

  const info = document.querySelector('#info');

  const freeTimingHelpers = [];

  function getTimingHelper() {
    return freeTimingHelpers.length > 0 ? freeTimingHelpers.pop() : new TimingHelper(device);
  }

  function freeTimingHelper(timingHelper) {
    freeTimingHelpers.push(timingHelper);
  }

  const code = `
  @group(0) @binding(0) var<storage,read_write> s: array<atomic<u32>>;

  @compute @workgroup_size(256, 1) fn cs(@builtin(global_invocation_id) id: vec3u) {
    let len = arrayLength(&s);
    atomicAdd(&s[len], 1);
  }
  `;

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

  const pipeline = device.createComputePipeline({
    layout: 'auto',
    compute: { module },
  });

  const range = (num, fn) => new Array(num).fill(0).map((_, i) => fn(i));
  const numBuffers = 32;
  const infos = range(numBuffers, () => {
    const storageBuffer = device.createBuffer({
      size: 1024 * 1024,
      usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST,
    });

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

    return {
      storageBuffer,
      bindGroup,
    };
  });

  const results = {
  };

  async function oneTimer() {
    const encoder = device.createCommandEncoder();
    const timingHelper = getTimingHelper();
    const pass = timingHelper.beginComputePass(encoder);
    pass.setPipeline(pipeline);

    infos.forEach(({ bindGroup }) => {
      pass.setBindGroup(0, bindGroup);
      pass.dispatchWorkgroups(32, 32);
    });

    pass.end();
    device.queue.submit([encoder.finish()]);

    const durationInNanoSeconds = await timingHelper.getResult();
    const durationInMilliseconds = durationInNanoSeconds / 1000 / 1000;
    freeTimingHelper(timingHelper);
    return durationInMilliseconds;
  }

  async function nTimers() {
    const encoder = device.createCommandEncoder();

    const timingHelpers = infos.map(({ bindGroup }) => {
      const timingHelper = getTimingHelper();
      const pass = timingHelper.beginComputePass(encoder);
      pass.setPipeline(pipeline);
      pass.setBindGroup(0, bindGroup);
      pass.dispatchWorkgroups(32, 32);
      pass.end();
      return timingHelper;
    });

    device.queue.submit([encoder.finish()]);

    const times = await Promise.all(timingHelpers.map(async (timingHelper) => {
      const durationInNanoSeconds = await timingHelper.getResult();
      const durationInMilliseconds = durationInNanoSeconds / 1000 / 1000;
      freeTimingHelper(timingHelper);
      return durationInMilliseconds;
    }));

    return times;
  }

  async function nEncoders() {
    const timingHelpers = infos.map(({ bindGroup }) => {
      const timingHelper = getTimingHelper();
      const encoder = device.createCommandEncoder();
      const pass = timingHelper.beginComputePass(encoder);
      pass.setPipeline(pipeline);
      pass.setBindGroup(0, bindGroup);
      pass.dispatchWorkgroups(32, 32);
      pass.end();
      device.queue.submit([encoder.finish()]);
      return timingHelper;
    });

    const times = await Promise.all(timingHelpers.map(async (timingHelper) => {
      const durationInNanoSeconds = await timingHelper.getResult();
      const durationInMilliseconds = durationInNanoSeconds / 1000 / 1000;
      freeTimingHelper(timingHelper);
      return durationInMilliseconds;
    }));

    return times;
  }

  let iter = 0;
  for (;;) {
    const oneTimerTime = await oneTimer();
    const nTimersTimes = await nTimers();
    const nEncoderTimes = await nEncoders();
    const nTimersTime = nTimersTimes.reduce((a, b) => a + b);
    const nEncoderTime = nTimersTimes.reduce((a, b) => a + b);
    info.textContent = `\
         iter: ${iter++}
 oneTimerTime: ${oneTimerTime}
  nTimersTime: ${nTimersTime}
nEncodersTime: ${nEncoderTime}

-- pass times --
    nTimer     nEncoder
${nTimersTimes.map((t, i) => `${i.toString().padStart(2)}: ${t.toFixed(4)}ms   ${nEncoderTimes[i].toFixed(4)}ms`).join('\n')}
`;
    await new Promise(resolve => setTimeout(resolve, 1000));
  }

}

function fail(msg) {
  const elem = document.querySelector('#fail');
  const contentElem = elem.querySelector('.content');
  elem.style.display = '';
  contentElem.textContent = msg;
}

main();
  

jsGist.json

{"name":"WebGPU Multiple Timestamps vs 1","settings":{},"filenames":["index.html","index.css","index.js"]}