Add packed integer dot product sample (#572)

Adds a small `dot4I8Packed` compute example toward #350.

The sample packs four signed 8-bit components into each `u32`, computes
four example dot products, and reads the results back from the GPU. A
dat.gui selector displays the input vectors and result for each example.

It requires `packed_4x8_integer_dot_product` and displays a requirement
message when that WGSL language feature is unavailable. There is one
packed compute pipeline, with no scalar fallback or in-sample result
validation.

### Testing

- Lint, source typecheck, and the Node 20.20.2 build pass. The build
emits existing typed-array warnings in unrelated samples.
- Chrome 153.0.8010.48, headless with `--enable-unsafe-webgpu`, on an
Apple Metal adapter: all four examples returned the expected results,
and the dat.gui selector was exercised.
- Desktop and narrow-screen layouts inspected. Injecting feature
unavailability displayed the requirement message without requesting an
adapter or compiling a shader.

No new dependencies or external assets.

---------

Co-authored-by: Kai Ninomiya <kainino@chromium.org>

Author: lyonsno

sample/packedIntegerDotProduct/index.html

@@ -0,0 +1,27 @@
+<!DOCTYPE html>
+<html lang="en">
+  <head>
+    <meta charset="utf-8" />
+    <meta name="viewport" content="width=device-width, initial-scale=1" />
+    <title>webgpu-samples: packedIntegerDotProduct</title>
+    <style>
+      :root {
+        color-scheme: light dark;
+      }
+      body {
+        width: calc(100vw - 280px);
+        height: 400px;
+        overflow-y: hidden;
+      }
+      pre {
+        white-space: pre-wrap;
+        height: 370px;
+      }
+    </style>
+    <script defer src="main.js" type="module"></script>
+    <script defer type="module" src="../../js/iframe-helper.js"></script>
+  </head>
+  <body>
+    <pre id="output">Computing packed integer dot products…</pre>
+  </body>
+</html>

sample/packedIntegerDotProduct/main.ts

@@ -0,0 +1,125 @@
+import { GUI } from 'dat.gui';
+import packedWGSL from './packed.wgsl';
+import { quitIfWebGPUNotAvailableOrMissingFeatures } from '../util';
+
+type vec4i = readonly [number, number, number, number];
+// Pack four signed 8-bit components into a u32, low byte first.
+function pack4xI8([x, y, z, w]: vec4i): number {
+  // `&` operator applies sign extension to i32 before operating.
+  // `>>> 0` converts the final i32 to u32.
+  /*prettier-ignore*/
+  return ((x & 0xff) |
+          ((y & 0xff) << 8) |
+          ((z & 0xff) << 16) |
+          ((w & 0xff) << 24)) >>> 0;
+}
+
+const outputElement = document.querySelector('#output') as HTMLElement;
+if (
+  !navigator.gpu?.wgslLanguageFeatures.has('packed_4x8_integer_dot_product')
+) {
+  result.textContent =
+    "This sample requires the WGSL language feature 'packed_4x8_integer_dot_product'.";
+} else {
+  const adapter = await navigator.gpu.requestAdapter({
+    featureLevel: 'compatibility',
+  });
+  const device = await adapter?.requestDevice();
+  quitIfWebGPUNotAvailableOrMissingFeatures(adapter, device);
+
+  const kInputSize = 2 * Uint32Array.BYTES_PER_ELEMENT;
+  const inputBuffer = device.createBuffer({
+    size: kInputSize,
+    usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.STORAGE,
+  });
+
+  const kOutputSize = Int32Array.BYTES_PER_ELEMENT;
+  const outputBuffer = device.createBuffer({
+    size: kOutputSize,
+    usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC,
+  });
+  const readbackBuffer = device.createBuffer({
+    size: kOutputSize,
+    usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ,
+  });
+
+  const pipeline = await device.createComputePipelineAsync({
+    layout: 'auto',
+    compute: { module: device.createShaderModule({ code: packedWGSL }) },
+  });
+  const bindGroup = device.createBindGroup({
+    layout: pipeline.getBindGroupLayout(0),
+    entries: [
+      { binding: 0, resource: { buffer: inputBuffer } },
+      { binding: 1, resource: { buffer: outputBuffer } },
+    ],
+  });
+
+  async function updateResult() {
+    // If an update is still in progress just wait until it's done.
+    if (readbackBuffer.mapState !== 'unmapped') {
+      setTimeout(updateResult, 0);
+      return;
+    }
+
+    const lhs = [settings.lhs0, settings.lhs1, settings.lhs2, settings.lhs3];
+    const rhs = [settings.rhs0, settings.rhs1, settings.rhs2, settings.rhs3];
+
+    device.queue.writeBuffer(
+      inputBuffer,
+      0,
+      new Uint32Array([lhs, rhs].map(pack4xI8))
+    );
+    const encoder = device.createCommandEncoder();
+    const pass = encoder.beginComputePass();
+    pass.setPipeline(pipeline);
+    pass.setBindGroup(0, bindGroup);
+    pass.dispatchWorkgroups(1);
+    pass.end();
+    encoder.copyBufferToBuffer(outputBuffer, 0, readbackBuffer, 0, kOutputSize);
+    device.queue.submit([encoder.finish()]);
+
+    await readbackBuffer.mapAsync(GPUMapMode.READ);
+    const result = new Int32Array(readbackBuffer.getMappedRange())[0];
+
+    // Result should be the same in JS, show that for comparison.
+    const expected =
+      lhs[0] * rhs[0] + lhs[1] * rhs[1] + lhs[2] * rhs[2] + lhs[3] * rhs[3];
+
+    const lhsStr = `[${lhs
+      .map((x) => x.toString().padStart(4))
+      .join(', ')}] (0x${pack4xI8(lhs).toString(16).padStart(8, '0')})`;
+    const rhsStr = `[${rhs
+      .map((x) => x.toString().padStart(4))
+      .join(', ')}] (0x${pack4xI8(rhs).toString(16).padStart(8, '0')})`;
+    const outStr = result.toString().padStart(6);
+    const expStr = expected.toString().padStart(6);
+    outputElement.textContent = `
+
+WGSL dot4I8Packed of ${lhsStr}
+                  by ${rhsStr} gave ${outStr} (JS gave ${expStr})`;
+
+    readbackBuffer.unmap();
+  }
+
+  const settings = {
+    lhs0: 1,
+    lhs1: -2,
+    lhs2: 3,
+    lhs3: -4,
+    rhs0: -5,
+    rhs1: 6,
+    rhs2: -7,
+    rhs3: 8,
+  };
+  const gui = new GUI();
+  gui.add(settings, 'lhs0', -127, 128, 1).onChange(updateResult);
+  gui.add(settings, 'lhs1', -127, 128, 1).onChange(updateResult);
+  gui.add(settings, 'lhs2', -127, 128, 1).onChange(updateResult);
+  gui.add(settings, 'lhs3', -127, 128, 1).onChange(updateResult);
+  gui.add(settings, 'rhs0', -127, 128, 1).onChange(updateResult);
+  gui.add(settings, 'rhs1', -127, 128, 1).onChange(updateResult);
+  gui.add(settings, 'rhs2', -127, 128, 1).onChange(updateResult);
+  gui.add(settings, 'rhs3', -127, 128, 1).onChange(updateResult);
+  updateResult();
+}

sample/packedIntegerDotProduct/meta.ts

@@ -0,0 +1,7 @@
+export default {
+  name: 'Packed Integer Dot Product',
+  description:
+    'Packs four signed 8-bit integers into each u32, computes their dot product with dot4I8Packed, and reads back the results.',
+  filename: __DIRNAME__,
+  sources: [{ path: 'main.ts' }, { path: 'packed.wgsl' }],
+};

sample/packedIntegerDotProduct/packed.wgsl

@@ -0,0 +1,11 @@
+requires packed_4x8_integer_dot_product;
+
+struct Input { lhs: u32, rhs: u32 }
+
+@group(0) @binding(0) var<storage, read> input: Input;
+@group(0) @binding(1) var<storage, read_write> output: i32;
+
+@compute @workgroup_size(1)
+fn main() {
+  output = dot4I8Packed(input.lhs, input.rhs);
+}

src/samples.ts

@@ -24,6 +24,7 @@ import normalMap from '../sample/normalMap/meta';
 import occlusionQuery from '../sample/occlusionQuery/meta';
 import particleLife from '../sample/particleLife/meta';
 import particles from '../sample/particles/meta';
+import packedIntegerDotProduct from '../sample/packedIntegerDotProduct/meta';
 import points from '../sample/points/meta';
 import primitivePicking from '../sample/primitivePicking/meta';
 import pristineGrid from '../sample/pristineGrid/meta';
@@ -116,6 +117,7 @@ export const pageCategories: PageCategory[] = [
       computeBoids,
       gameOfLife,
       bitonicSort,
+      packedIntegerDotProduct,
     },
   },