scan-negative-zero.html

<!DOCTYPE html>

<html>
  <head>
    <meta charset="utf-8" />
    <title>Tiny Exclusive Scan Test</title>
  </head>

  <body>
    <script type="module">
      // begin TimingHelper code
      function assert(cond, msg = "") {
        if (!cond) {
          throw new Error(msg);
        }
      }

      const adapter = await navigator.gpu?.requestAdapter();
      const canTimestamp = adapter.features.has("timestamp-query");
      const hasSubgroups = adapter.features.has("subgroups");
      const device = await adapter?.requestDevice({
        requiredFeatures: [
          ...(canTimestamp ? ["timestamp-query"] : []),
          ...(hasSubgroups ? ["subgroups"] : []),
        ],
      });

      if (!device) {
        fail("Fatal error: Device does not support WebGPU.");
      }

      if (!canTimestamp) {
        fail(
          'Fatal error: Device does not support WebGPU timestamp query (`adapter.features.has("timestamp-query")` is false).'
        );
      }

      const workgroupSize = 64;
      const memsrcSize = 64;

      const workgroupCount = memsrcSize / workgroupSize;
      const dispatchGeometry = [workgroupCount, 1];
      while (
        dispatchGeometry[0] > adapter.limits.maxComputeWorkgroupsPerDimension
      ) {
        dispatchGeometry[0] /= 2;
        dispatchGeometry[1] *= 2;
      }
      console.log(`workgroup count: ${workgroupCount}
      workgroup size: ${workgroupSize}
      maxComputeWGPerDim: ${adapter.limits.maxComputeWorkgroupsPerDimension}
      dispatchGeometry: ${dispatchGeometry}`);

      const memsrc = new Float32Array(memsrcSize);
      const ref = new Float32Array(memsrcSize);
      for (let i = 0; i < memsrc.length; i++) {
        memsrc[i] = i;
        ref[i] = i % 32 == 0 ? 0 : i - 1 + ref[i - 1];
      }

      const memcpyModule = device.createShaderModule({
        label: "exclusive-scan subgroups kernel",
        code: /* wgsl */ `
        enable subgroups;
        /* output */
        @group(0) @binding(0) var<storage, read_write> memDest: array<f32>;
        /* input */
        @group(0) @binding(1) var<storage, read> memSrc: array<f32>;

        @compute @workgroup_size(${workgroupSize}) fn memcpyKernel(
          @builtin(global_invocation_id) id: vec3u,
          @builtin(num_workgroups) nwg: vec3u,
          @builtin(workgroup_id) wgid: vec3u) {
          let i = id.y * nwg.x * ${workgroupSize} + id.x;
          let in = memSrc[i];
          memDest[i] = subgroupExclusiveAdd(in);
        }`,
      });

      const memcpyPipeline = device.createComputePipeline({
        label: "memcpy compute pipeline",
        layout: "auto",
        compute: {
          module: memcpyModule,
        },
      });

      // create buffers on the GPU to hold data
      // read-only inputs:
      const memsrcBuffer = device.createBuffer({
        label: "memory source buffer",
        size: memsrc.byteLength,
        usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST,
      });
      device.queue.writeBuffer(memsrcBuffer, 0, memsrc);

      const memdestBuffer = device.createBuffer({
        label: "memory destination buffer",
        size: memsrc.byteLength,
        usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC,
      });

      const mappableMemdstBuffer = device.createBuffer({
        label: "mappable memory destination buffer",
        size: memsrc.byteLength,
        usage: GPUBufferUsage.MAP_READ | GPUBufferUsage.COPY_DST,
      });

      /** Set up bindGroups per compute kernel to tell the shader which buffers to use */
      const memcpyBindGroup = device.createBindGroup({
        label: "bindGroup for memcpy kernel",
        layout: memcpyPipeline.getBindGroupLayout(0),
        entries: [
          { binding: 0, resource: { buffer: memdestBuffer } },
          { binding: 1, resource: { buffer: memsrcBuffer } },
        ],
      });

      const encoder = device.createCommandEncoder({
        label: "memcpy encoder",
      });

      const memcpyPass = encoder.beginComputePass(encoder, {
        label: "memcpy compute pass",
      });
      memcpyPass.setPipeline(memcpyPipeline);
      memcpyPass.setBindGroup(0, memcpyBindGroup);
      // TODO handle not evenly divisible by wgSize
      memcpyPass.dispatchWorkgroups(...dispatchGeometry);
      memcpyPass.end();

      // Encode a command to copy the results to a mappable buffer.
      // this is (from, to)
      encoder.copyBufferToBuffer(
        memdestBuffer,
        0,
        mappableMemdstBuffer,
        0,
        mappableMemdstBuffer.size
      );

      // Finish encoding and submit the commands
      const command_buffer = encoder.finish();
      device.queue.submit([command_buffer]);

      // Read the results
      await mappableMemdstBuffer.mapAsync(GPUMapMode.READ);
      const memdest = new Float32Array(
        mappableMemdstBuffer.getMappedRange().slice()
      );
      mappableMemdstBuffer.unmap();
      let errors = 0;
      for (let i = 0; i < memdest.length; i++) {
        if (memdest[i] != ref[i]) {
          if (errors < 5) {
            console.log(
              `Error ${errors}: i=${i}, src=${memsrc[i]}, dest=${memdest[i]}, expected=${ref[i]}`
            );
          }
          errors++;
        }
      }
      if (errors > 0) {
        console.log(`Memdest size: ${memdest.length} | Errors: ${errors}`);
      } else {
        console.log(`Memdest size: ${memdest.length} | No errors!`);
      }
      console.log("Input", memsrc, "Output", memdest);

      function fail(msg) {
        // eslint-disable-next-line no-alert
        alert(msg);
      }
    </script>
  </body>
</html>