From 2e6bba3ebb75c201b4d39f9cc2e52de3c80ee39e Mon Sep 17 00:00:00 2001 From: Noah Lyons Date: Sat, 11 Jul 2026 21:48:58 -0400 Subject: [PATCH 1/6] Add packed integer dot product sample --- package.json | 1 + sample/packedIntegerDotProduct/index.html | 151 +++++++++++++++ sample/packedIntegerDotProduct/main.ts | 203 ++++++++++++++++++++ sample/packedIntegerDotProduct/meta.ts | 12 ++ sample/packedIntegerDotProduct/packed.wgsl | 13 ++ sample/packedIntegerDotProduct/reference.ts | 117 +++++++++++ sample/packedIntegerDotProduct/scalar.wgsl | 24 +++ src/samples.ts | 2 + test/packedIntegerDotProduct.test.mjs | 91 +++++++++ 9 files changed, 614 insertions(+) create mode 100644 sample/packedIntegerDotProduct/index.html create mode 100644 sample/packedIntegerDotProduct/main.ts create mode 100644 sample/packedIntegerDotProduct/meta.ts create mode 100644 sample/packedIntegerDotProduct/packed.wgsl create mode 100644 sample/packedIntegerDotProduct/reference.ts create mode 100644 sample/packedIntegerDotProduct/scalar.wgsl create mode 100644 test/packedIntegerDotProduct.test.mjs diff --git a/package.json b/package.json index 72d88853..1c50f22c 100644 --- a/package.json +++ b/package.json @@ -10,6 +10,7 @@ "url": "https://github.com/webgpu/webgpu-samples.git" }, "scripts": { + "test": "node --test test/packedIntegerDotProduct.test.mjs", "lint": "eslint --ext .ts,.js,.html src/ sample/ build/ .eslintrc.cjs rollup.config.js index.html", "fix": "eslint --fix --ext .ts,.js,.html src/ sample/ build/ .eslintrc.cjs rollup.config.js index.html", "build": "node build/tools/build.js", diff --git a/sample/packedIntegerDotProduct/index.html b/sample/packedIntegerDotProduct/index.html new file mode 100644 index 00000000..c62baa66 --- /dev/null +++ b/sample/packedIntegerDotProduct/index.html @@ -0,0 +1,151 @@ + + + + + + webgpu-samples: packedIntegerDotProduct + + + + + +
+
+

Packed Integer Dot Product

+
+ + + +
+
+ +
+
+
Effective route
+
Pending
+
+
+
Packed support
+
Checking
+
+
+
Readback time
+
Pending
+
+
+
Validation
+
Pending
+
+
+

+
+ + diff --git a/sample/packedIntegerDotProduct/main.ts b/sample/packedIntegerDotProduct/main.ts new file mode 100644 index 00000000..092733d4 --- /dev/null +++ b/sample/packedIntegerDotProduct/main.ts @@ -0,0 +1,203 @@ +import packedWGSL from './packed.wgsl'; +import scalarWGSL from './scalar.wgsl'; +import { + makeInputVectors, + packedDotLanguageFeature, + RequestedRoute, + selectRoute, + validateRun, +} from './reference'; +import { quitIfWebGPUNotAvailableOrMissingFeatures } from '../util'; + +const elementCount = 64 * 1024; +const workgroupSize = 64; +const previewCount = 192; +const { lhsPacked, rhsPacked, expected } = makeInputVectors(elementCount); + +const adapter = await navigator.gpu?.requestAdapter({ + featureLevel: 'compatibility', +}); +const device = await adapter?.requestDevice(); +quitIfWebGPUNotAvailableOrMissingFeatures(adapter, device); + +const packedLanguageFeatureSupported = navigator.gpu.wgslLanguageFeatures.has( + packedDotLanguageFeature +); +const languageFeatures = new Set( + packedLanguageFeatureSupported ? [packedDotLanguageFeature] : [] +); + +function createInputBuffer(data: Uint32Array) { + const buffer = device.createBuffer({ + size: data.byteLength, + usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST, + }); + device.queue.writeBuffer( + buffer, + 0, + data.buffer as ArrayBuffer, + data.byteOffset, + data.byteLength + ); + return buffer; +} + +const lhsBuffer = createInputBuffer(lhsPacked); +const rhsBuffer = createInputBuffer(rhsPacked); +const outputBuffer = device.createBuffer({ + size: expected.byteLength, + usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC, +}); +const readbackBuffer = device.createBuffer({ + size: expected.byteLength, + usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ, +}); + +async function createPipeline(code: string) { + return device.createComputePipelineAsync({ + layout: 'auto', + compute: { + module: device.createShaderModule({ code }), + }, + }); +} + +const scalarPipeline = await createPipeline(scalarWGSL); +const packedPipeline = packedLanguageFeatureSupported + ? await createPipeline(packedWGSL) + : null; + +function createBindGroup(pipeline: GPUComputePipeline) { + return device.createBindGroup({ + layout: pipeline.getBindGroupLayout(0), + entries: [ + { binding: 0, resource: { buffer: lhsBuffer } }, + { binding: 1, resource: { buffer: rhsBuffer } }, + { binding: 2, resource: { buffer: outputBuffer } }, + ], + }); +} + +const scalarBindGroup = createBindGroup(scalarPipeline); +const packedBindGroup = packedPipeline ? createBindGroup(packedPipeline) : null; + +const routeValue = document.querySelector('[data-route-value]') as HTMLElement; +const supportValue = document.querySelector( + '[data-support-value]' +) as HTMLElement; +const durationValue = document.querySelector( + '[data-duration-value]' +) as HTMLElement; +const validationValue = document.querySelector( + '[data-validation-value]' +) as HTMLElement; +const fallbackValue = document.querySelector( + '[data-fallback-value]' +) as HTMLElement; +const canvas = document.querySelector('canvas') as HTMLCanvasElement; +const context = canvas.getContext('2d') as CanvasRenderingContext2D; +const buttons = Array.from( + document.querySelectorAll('[data-route]') +); + +supportValue.textContent = packedLanguageFeatureSupported + ? 'Available' + : 'Unavailable'; + +function drawResults(results: Int32Array) { + const scale = window.devicePixelRatio; + const width = canvas.clientWidth; + const height = canvas.clientHeight; + canvas.width = Math.round(width * scale); + canvas.height = Math.round(height * scale); + context.setTransform(scale, 0, 0, scale, 0, 0); + context.clearRect(0, 0, width, height); + + const baseline = height / 2; + const maxMagnitude = Math.max( + 1, + ...results.slice(0, previewCount).map((value) => Math.abs(value)) + ); + const barWidth = width / previewCount; + + context.fillStyle = '#d8dde5'; + context.fillRect(0, baseline, width, 1); + for (let index = 0; index < previewCount; ++index) { + const value = results[index]; + const barHeight = (Math.abs(value) / maxMagnitude) * (baseline - 12); + context.fillStyle = value >= 0 ? '#00a87a' : '#e34b7a'; + context.fillRect( + index * barWidth, + value >= 0 ? baseline - barHeight : baseline + 1, + Math.max(1, barWidth - 1), + barHeight + ); + } +} + +let runSerial = 0; + +async function run(requestedRoute: RequestedRoute) { + const serial = ++runSerial; + buttons.forEach((button) => { + button.disabled = true; + button.dataset.selected = String(button.dataset.route === requestedRoute); + }); + validationValue.textContent = 'Running'; + fallbackValue.textContent = ''; + + const selection = selectRoute(requestedRoute, languageFeatures); + const pipeline = + selection.effectiveRoute === 'packed' ? packedPipeline : scalarPipeline; + const bindGroup = + selection.effectiveRoute === 'packed' ? packedBindGroup : scalarBindGroup; + if (!pipeline || !bindGroup) { + throw new Error( + `effective route '${selection.effectiveRoute}' is unavailable` + ); + } + + const encoder = device.createCommandEncoder(); + const pass = encoder.beginComputePass(); + pass.setPipeline(pipeline); + pass.setBindGroup(0, bindGroup); + pass.dispatchWorkgroups(Math.ceil(elementCount / workgroupSize)); + pass.end(); + encoder.copyBufferToBuffer( + outputBuffer, + 0, + readbackBuffer, + 0, + expected.byteLength + ); + + const start = performance.now(); + device.queue.submit([encoder.finish()]); + await readbackBuffer.mapAsync(GPUMapMode.READ); + const duration = performance.now() - start; + const actual = new Int32Array(readbackBuffer.getMappedRange()).slice(); + readbackBuffer.unmap(); + + validateRun({ ...selection, expected, actual }); + if (serial !== runSerial) { + return; + } + + routeValue.textContent = selection.effectiveRoute; + durationValue.textContent = `${duration.toFixed(2)} ms`; + validationValue.textContent = `${actual.length.toLocaleString()} exact`; + fallbackValue.textContent = selection.fallbackReason ?? ''; + drawResults(actual); + buttons.forEach((button) => { + button.disabled = false; + }); +} + +buttons.forEach((button) => { + button.addEventListener('click', () => { + run(button.dataset.route as RequestedRoute); + }); +}); + +window.addEventListener('resize', () => drawResults(expected)); +await run('auto'); diff --git a/sample/packedIntegerDotProduct/meta.ts b/sample/packedIntegerDotProduct/meta.ts new file mode 100644 index 00000000..d213e688 --- /dev/null +++ b/sample/packedIntegerDotProduct/meta.ts @@ -0,0 +1,12 @@ +export default { + name: 'Packed Integer Dot Product', + description: + 'Compares signed 8-bit dot products using the packed WGSL instruction and an equivalent scalar route, with exact result validation.', + filename: __DIRNAME__, + sources: [ + { path: 'main.ts' }, + { path: 'reference.ts' }, + { path: 'packed.wgsl' }, + { path: 'scalar.wgsl' }, + ], +}; diff --git a/sample/packedIntegerDotProduct/packed.wgsl b/sample/packedIntegerDotProduct/packed.wgsl new file mode 100644 index 00000000..46b68fde --- /dev/null +++ b/sample/packedIntegerDotProduct/packed.wgsl @@ -0,0 +1,13 @@ +requires packed_4x8_integer_dot_product; + +@group(0) @binding(0) var lhs: array; +@group(0) @binding(1) var rhs: array; +@group(0) @binding(2) var output: array; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) id: vec3) { + if (id.x >= arrayLength(&output)) { + return; + } + output[id.x] = dot4I8Packed(lhs[id.x], rhs[id.x]); +} diff --git a/sample/packedIntegerDotProduct/reference.ts b/sample/packedIntegerDotProduct/reference.ts new file mode 100644 index 00000000..26a70b1a --- /dev/null +++ b/sample/packedIntegerDotProduct/reference.ts @@ -0,0 +1,117 @@ +export const packedDotLanguageFeature = 'packed_4x8_integer_dot_product'; + +export type RequestedRoute = 'auto' | 'packed' | 'scalar'; +export type EffectiveRoute = 'packed' | 'scalar'; + +export type RouteSelection = { + requestedRoute: RequestedRoute; + effectiveRoute: EffectiveRoute; + packedLanguageFeatureSupported: boolean; + fallbackReason: string | null; +}; + +function assertI8(value: number) { + if (!Number.isInteger(value) || value < -128 || value > 127) { + throw new RangeError(`expected signed 8-bit integer, received ${value}`); + } +} + +export function pack4xI8(values: readonly number[]): number { + if (values.length !== 4) { + throw new RangeError(`expected 4 values, received ${values.length}`); + } + + let packed = 0; + values.forEach((value, index) => { + assertI8(value); + packed |= (value & 0xff) << (index * 8); + }); + return packed >>> 0; +} + +export function dot4I8(lhs: readonly number[], rhs: readonly number[]): number { + if (lhs.length !== 4 || rhs.length !== 4) { + throw new RangeError('signed dot product requires two 4-component vectors'); + } + + let result = 0; + for (let index = 0; index < 4; ++index) { + assertI8(lhs[index]); + assertI8(rhs[index]); + result += lhs[index] * rhs[index]; + } + return result; +} + +export function selectRoute( + requestedRoute: RequestedRoute, + languageFeatures: ReadonlySet +): RouteSelection { + const packedLanguageFeatureSupported = languageFeatures.has( + packedDotLanguageFeature + ); + const wantsPacked = requestedRoute !== 'scalar'; + const effectiveRoute = + wantsPacked && packedLanguageFeatureSupported ? 'packed' : 'scalar'; + + return { + requestedRoute, + effectiveRoute, + packedLanguageFeatureSupported, + fallbackReason: + requestedRoute === 'packed' && !packedLanguageFeatureSupported + ? `WGSL language feature '${packedDotLanguageFeature}' is unavailable` + : null, + }; +} + +export function validateRun({ + requestedRoute, + effectiveRoute, + packedLanguageFeatureSupported, + expected, + actual, +}: RouteSelection & { + expected: ArrayLike; + actual: ArrayLike; +}) { + if (effectiveRoute === 'packed' && !packedLanguageFeatureSupported) { + throw new Error( + `requested route '${requestedRoute}' claims packed route without language feature support` + ); + } + if (actual.length !== expected.length) { + throw new Error( + `partial output: expected ${expected.length} results, received ${actual.length}` + ); + } + for (let index = 0; index < expected.length; ++index) { + if (actual[index] !== expected[index]) { + throw new Error( + `result ${index} mismatch: expected ${expected[index]}, received ${actual[index]}` + ); + } + } +} + +export function makeInputVectors(count: number) { + const lhsPacked = new Uint32Array(count); + const rhsPacked = new Uint32Array(count); + const expected = new Int32Array(count); + let state = 0x13579bdf; + + const nextI8 = () => { + state = (Math.imul(state, 1664525) + 1013904223) >>> 0; + return ((state >>> 24) & 0xff) - 128; + }; + + for (let index = 0; index < count; ++index) { + const lhs = [nextI8(), nextI8(), nextI8(), nextI8()]; + const rhs = [nextI8(), nextI8(), nextI8(), nextI8()]; + lhsPacked[index] = pack4xI8(lhs); + rhsPacked[index] = pack4xI8(rhs); + expected[index] = dot4I8(lhs, rhs); + } + + return { lhsPacked, rhsPacked, expected }; +} diff --git a/sample/packedIntegerDotProduct/scalar.wgsl b/sample/packedIntegerDotProduct/scalar.wgsl new file mode 100644 index 00000000..7bf067e4 --- /dev/null +++ b/sample/packedIntegerDotProduct/scalar.wgsl @@ -0,0 +1,24 @@ +@group(0) @binding(0) var lhs: array; +@group(0) @binding(1) var rhs: array; +@group(0) @binding(2) var output: array; + +fn unpackI8(value: u32, component: u32) -> i32 { + let byte = i32((value >> (component * 8u)) & 0xffu); + return select(byte, byte - 256, byte >= 128); +} + +fn dot4I8Scalar(a: u32, b: u32) -> i32 { + var result = 0; + for (var component = 0u; component < 4u; component += 1u) { + result += unpackI8(a, component) * unpackI8(b, component); + } + return result; +} + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) id: vec3) { + if (id.x >= arrayLength(&output)) { + return; + } + output[id.x] = dot4I8Scalar(lhs[id.x], rhs[id.x]); +} diff --git a/src/samples.ts b/src/samples.ts index c50d6c45..a807048f 100644 --- a/src/samples.ts +++ b/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, }, }, diff --git a/test/packedIntegerDotProduct.test.mjs b/test/packedIntegerDotProduct.test.mjs new file mode 100644 index 00000000..5a952851 --- /dev/null +++ b/test/packedIntegerDotProduct.test.mjs @@ -0,0 +1,91 @@ +import assert from 'node:assert/strict'; +import test from 'node:test'; + +const referenceModuleUrl = new URL( + '../sample/packedIntegerDotProduct/reference.ts', + import.meta.url +); + +async function loadReferenceModule() { + try { + return await import(referenceModuleUrl); + } catch (error) { + if (error?.code === 'ERR_MODULE_NOT_FOUND') { + assert.fail( + 'packed integer dot-product reference and route contract is missing' + ); + } + throw error; + } +} + +test('signed packed dot products match scalar arithmetic', async () => { + const { dot4I8, pack4xI8 } = await loadReferenceModule(); + const lhs = [1, -2, 3, -4]; + const rhs = [-5, 6, -7, 8]; + + assert.equal(pack4xI8([1, -2, 127, -128]), 0x807ffe01); + assert.equal(dot4I8(lhs, rhs), -70); +}); + +test('route selection records requested and effective identity', async () => { + const { selectRoute } = await loadReferenceModule(); + const feature = 'packed_4x8_integer_dot_product'; + + assert.deepEqual(selectRoute('auto', new Set([feature])), { + requestedRoute: 'auto', + effectiveRoute: 'packed', + packedLanguageFeatureSupported: true, + fallbackReason: null, + }); + assert.deepEqual(selectRoute('packed', new Set()), { + requestedRoute: 'packed', + effectiveRoute: 'scalar', + packedLanguageFeatureSupported: false, + fallbackReason: `WGSL language feature '${feature}' is unavailable`, + }); + assert.deepEqual(selectRoute('scalar', new Set([feature])), { + requestedRoute: 'scalar', + effectiveRoute: 'scalar', + packedLanguageFeatureSupported: true, + fallbackReason: null, + }); +}); + +test('evidence validation rejects wrong routes and partial output', async () => { + const { validateRun } = await loadReferenceModule(); + + assert.throws( + () => + validateRun({ + requestedRoute: 'packed', + effectiveRoute: 'packed', + packedLanguageFeatureSupported: false, + expected: [11, 22], + actual: [11, 22], + }), + /claims packed route without language feature support/ + ); + assert.throws( + () => + validateRun({ + requestedRoute: 'auto', + effectiveRoute: 'scalar', + packedLanguageFeatureSupported: false, + expected: [11, 22], + actual: [11], + }), + /partial output: expected 2 results, received 1/ + ); + assert.throws( + () => + validateRun({ + requestedRoute: 'auto', + effectiveRoute: 'scalar', + packedLanguageFeatureSupported: false, + expected: [11, 22], + actual: [11, 23], + }), + /result 1 mismatch: expected 22, received 23/ + ); +}); From 4b7624bc644ea6cc31856a09ca18f2aca990df59 Mon Sep 17 00:00:00 2001 From: Noah Lyons Date: Tue, 15 Sep 2026 21:35:14 -0400 Subject: [PATCH 2/6] Compile the packed dot product reference for Node 20 tests --- test/packedIntegerDotProduct.test.mjs | 22 ++++++++++++---------- 1 file changed, 12 insertions(+), 10 deletions(-) diff --git a/test/packedIntegerDotProduct.test.mjs b/test/packedIntegerDotProduct.test.mjs index 5a952851..59376e54 100644 --- a/test/packedIntegerDotProduct.test.mjs +++ b/test/packedIntegerDotProduct.test.mjs @@ -1,5 +1,7 @@ import assert from 'node:assert/strict'; +import { readFile } from 'node:fs/promises'; import test from 'node:test'; +import ts from 'typescript'; const referenceModuleUrl = new URL( '../sample/packedIntegerDotProduct/reference.ts', @@ -7,16 +9,16 @@ const referenceModuleUrl = new URL( ); async function loadReferenceModule() { - try { - return await import(referenceModuleUrl); - } catch (error) { - if (error?.code === 'ERR_MODULE_NOT_FOUND') { - assert.fail( - 'packed integer dot-product reference and route contract is missing' - ); - } - throw error; - } + // Node 20 cannot import TypeScript directly. Compile this standalone module + // with the same TypeScript dependency used by the sample build. + const source = await readFile(referenceModuleUrl, 'utf8'); + const { outputText } = ts.transpileModule(source, { + compilerOptions: { + target: ts.ScriptTarget.ES2022, + module: ts.ModuleKind.ESNext, + }, + }); + return import(`data:text/javascript,${encodeURIComponent(outputText)}`); } test('signed packed dot products match scalar arithmetic', async () => { From bd3a0f1b5d634b6e15a19ef644b407994a8611bf Mon Sep 17 00:00:00 2001 From: Noah Lyons Date: Tue, 15 Sep 2026 23:50:42 -0400 Subject: [PATCH 3/6] Simplify packed integer dot product sample for learners --- package.json | 1 - sample/packedIntegerDotProduct/index.html | 132 +--------- sample/packedIntegerDotProduct/main.ts | 260 ++++++-------------- sample/packedIntegerDotProduct/meta.ts | 9 +- sample/packedIntegerDotProduct/reference.ts | 117 --------- sample/packedIntegerDotProduct/scalar.wgsl | 24 -- test/packedIntegerDotProduct.test.mjs | 93 ------- 7 files changed, 88 insertions(+), 548 deletions(-) delete mode 100644 sample/packedIntegerDotProduct/reference.ts delete mode 100644 sample/packedIntegerDotProduct/scalar.wgsl delete mode 100644 test/packedIntegerDotProduct.test.mjs diff --git a/package.json b/package.json index 1c50f22c..72d88853 100644 --- a/package.json +++ b/package.json @@ -10,7 +10,6 @@ "url": "https://github.com/webgpu/webgpu-samples.git" }, "scripts": { - "test": "node --test test/packedIntegerDotProduct.test.mjs", "lint": "eslint --ext .ts,.js,.html src/ sample/ build/ .eslintrc.cjs rollup.config.js index.html", "fix": "eslint --fix --ext .ts,.js,.html src/ sample/ build/ .eslintrc.cjs rollup.config.js index.html", "build": "node build/tools/build.js", diff --git a/sample/packedIntegerDotProduct/index.html b/sample/packedIntegerDotProduct/index.html index c62baa66..615321d8 100644 --- a/sample/packedIntegerDotProduct/index.html +++ b/sample/packedIntegerDotProduct/index.html @@ -7,145 +7,23 @@ -
-
-

Packed Integer Dot Product

-
- - - -
-
- -
-
-
Effective route
-
Pending
-
-
-
Packed support
-
Checking
-
-
-
Readback time
-
Pending
-
-
-
Validation
-
Pending
-
-
-

-
+ 
Computing packed integer dot products…
diff --git a/sample/packedIntegerDotProduct/main.ts b/sample/packedIntegerDotProduct/main.ts index 092733d4..007a71f3 100644 --- a/sample/packedIntegerDotProduct/main.ts +++ b/sample/packedIntegerDotProduct/main.ts @@ -1,203 +1,105 @@ +import { GUI } from 'dat.gui'; import packedWGSL from './packed.wgsl'; -import scalarWGSL from './scalar.wgsl'; -import { - makeInputVectors, - packedDotLanguageFeature, - RequestedRoute, - selectRoute, - validateRun, -} from './reference'; import { quitIfWebGPUNotAvailableOrMissingFeatures } from '../util'; -const elementCount = 64 * 1024; -const workgroupSize = 64; -const previewCount = 192; -const { lhsPacked, rhsPacked, expected } = makeInputVectors(elementCount); - -const adapter = await navigator.gpu?.requestAdapter({ - featureLevel: 'compatibility', -}); -const device = await adapter?.requestDevice(); -quitIfWebGPUNotAvailableOrMissingFeatures(adapter, device); - -const packedLanguageFeatureSupported = navigator.gpu.wgslLanguageFeatures.has( - packedDotLanguageFeature -); -const languageFeatures = new Set( - packedLanguageFeatureSupported ? [packedDotLanguageFeature] : [] -); - -function createInputBuffer(data: Uint32Array) { - const buffer = device.createBuffer({ - size: data.byteLength, - usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST, +const result = document.querySelector('#result') 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', }); - device.queue.writeBuffer( - buffer, - 0, - data.buffer as ArrayBuffer, - data.byteOffset, - data.byteLength - ); - return buffer; -} - -const lhsBuffer = createInputBuffer(lhsPacked); -const rhsBuffer = createInputBuffer(rhsPacked); -const outputBuffer = device.createBuffer({ - size: expected.byteLength, - usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC, -}); -const readbackBuffer = device.createBuffer({ - size: expected.byteLength, - usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ, -}); + const device = await adapter?.requestDevice(); + quitIfWebGPUNotAvailableOrMissingFeatures(adapter, device); + + const lhs = [ + [1, -2, 3, -4], + [-128, -128, -128, -128], + [127, 127, 127, 127], + [1, 2, 3, 4], + ]; + const rhs = [ + [-5, 6, -7, 8], + [-128, -128, -128, -128], + [127, 127, 127, 127], + [4, 3, 2, 1], + ]; + + function createInputBuffer(vectors: number[][]) { + // Pack four signed 8-bit components into each u32, low byte first. + // Masking preserves the two's-complement representation of negative values. + const packed = new Uint32Array( + vectors.map( + ([x, y, z, w]) => + (x & 0xff) | + ((y & 0xff) << 8) | + ((z & 0xff) << 16) | + ((w & 0xff) << 24) + ) + ); + const buffer = device.createBuffer({ + size: packed.byteLength, + usage: GPUBufferUsage.STORAGE, + mappedAtCreation: true, + }); + new Uint32Array(buffer.getMappedRange()).set(packed); + buffer.unmap(); + return buffer; + } -async function createPipeline(code: string) { - return device.createComputePipelineAsync({ + const outputSize = lhs.length * Int32Array.BYTES_PER_ELEMENT; + const outputBuffer = device.createBuffer({ + size: outputSize, + usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC, + }); + const readbackBuffer = device.createBuffer({ + size: outputSize, + usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ, + }); + const pipeline = await device.createComputePipelineAsync({ layout: 'auto', - compute: { - module: device.createShaderModule({ code }), - }, + compute: { module: device.createShaderModule({ code: packedWGSL }) }, }); -} - -const scalarPipeline = await createPipeline(scalarWGSL); -const packedPipeline = packedLanguageFeatureSupported - ? await createPipeline(packedWGSL) - : null; - -function createBindGroup(pipeline: GPUComputePipeline) { - return device.createBindGroup({ + const bindGroup = device.createBindGroup({ layout: pipeline.getBindGroupLayout(0), entries: [ - { binding: 0, resource: { buffer: lhsBuffer } }, - { binding: 1, resource: { buffer: rhsBuffer } }, + { binding: 0, resource: { buffer: createInputBuffer(lhs) } }, + { binding: 1, resource: { buffer: createInputBuffer(rhs) } }, { binding: 2, resource: { buffer: outputBuffer } }, ], }); -} - -const scalarBindGroup = createBindGroup(scalarPipeline); -const packedBindGroup = packedPipeline ? createBindGroup(packedPipeline) : null; - -const routeValue = document.querySelector('[data-route-value]') as HTMLElement; -const supportValue = document.querySelector( - '[data-support-value]' -) as HTMLElement; -const durationValue = document.querySelector( - '[data-duration-value]' -) as HTMLElement; -const validationValue = document.querySelector( - '[data-validation-value]' -) as HTMLElement; -const fallbackValue = document.querySelector( - '[data-fallback-value]' -) as HTMLElement; -const canvas = document.querySelector('canvas') as HTMLCanvasElement; -const context = canvas.getContext('2d') as CanvasRenderingContext2D; -const buttons = Array.from( - document.querySelectorAll('[data-route]') -); - -supportValue.textContent = packedLanguageFeatureSupported - ? 'Available' - : 'Unavailable'; - -function drawResults(results: Int32Array) { - const scale = window.devicePixelRatio; - const width = canvas.clientWidth; - const height = canvas.clientHeight; - canvas.width = Math.round(width * scale); - canvas.height = Math.round(height * scale); - context.setTransform(scale, 0, 0, scale, 0, 0); - context.clearRect(0, 0, width, height); - - const baseline = height / 2; - const maxMagnitude = Math.max( - 1, - ...results.slice(0, previewCount).map((value) => Math.abs(value)) - ); - const barWidth = width / previewCount; - - context.fillStyle = '#d8dde5'; - context.fillRect(0, baseline, width, 1); - for (let index = 0; index < previewCount; ++index) { - const value = results[index]; - const barHeight = (Math.abs(value) / maxMagnitude) * (baseline - 12); - context.fillStyle = value >= 0 ? '#00a87a' : '#e34b7a'; - context.fillRect( - index * barWidth, - value >= 0 ? baseline - barHeight : baseline + 1, - Math.max(1, barWidth - 1), - barHeight - ); - } -} - -let runSerial = 0; - -async function run(requestedRoute: RequestedRoute) { - const serial = ++runSerial; - buttons.forEach((button) => { - button.disabled = true; - button.dataset.selected = String(button.dataset.route === requestedRoute); - }); - validationValue.textContent = 'Running'; - fallbackValue.textContent = ''; - - const selection = selectRoute(requestedRoute, languageFeatures); - const pipeline = - selection.effectiveRoute === 'packed' ? packedPipeline : scalarPipeline; - const bindGroup = - selection.effectiveRoute === 'packed' ? packedBindGroup : scalarBindGroup; - if (!pipeline || !bindGroup) { - throw new Error( - `effective route '${selection.effectiveRoute}' is unavailable` - ); - } const encoder = device.createCommandEncoder(); const pass = encoder.beginComputePass(); pass.setPipeline(pipeline); pass.setBindGroup(0, bindGroup); - pass.dispatchWorkgroups(Math.ceil(elementCount / workgroupSize)); + pass.dispatchWorkgroups(Math.ceil(lhs.length / 64)); pass.end(); - encoder.copyBufferToBuffer( - outputBuffer, - 0, - readbackBuffer, - 0, - expected.byteLength - ); - - const start = performance.now(); + encoder.copyBufferToBuffer(outputBuffer, 0, readbackBuffer, 0, outputSize); device.queue.submit([encoder.finish()]); + await readbackBuffer.mapAsync(GPUMapMode.READ); - const duration = performance.now() - start; - const actual = new Int32Array(readbackBuffer.getMappedRange()).slice(); + const results = new Int32Array(readbackBuffer.getMappedRange()).slice(); readbackBuffer.unmap(); - validateRun({ ...selection, expected, actual }); - if (serial !== runSerial) { - return; + const settings = { example: 0 }; + function showResult() { + const i = settings.example; + result.textContent = `a = [${lhs[i].join(', ')}] +b = [${rhs[i].join(', ')}] +dot4I8Packed(a, b) = ${results[i]}`; } - - routeValue.textContent = selection.effectiveRoute; - durationValue.textContent = `${duration.toFixed(2)} ms`; - validationValue.textContent = `${actual.length.toLocaleString()} exact`; - fallbackValue.textContent = selection.fallbackReason ?? ''; - drawResults(actual); - buttons.forEach((button) => { - button.disabled = false; - }); + const gui = new GUI(); + gui + .add(settings, 'example', { + 'Mixed signs': 0, + 'Minimum signed bytes': 1, + 'Maximum signed bytes': 2, + 'Positive components': 3, + }) + .onChange(showResult); + showResult(); } - -buttons.forEach((button) => { - button.addEventListener('click', () => { - run(button.dataset.route as RequestedRoute); - }); -}); - -window.addEventListener('resize', () => drawResults(expected)); -await run('auto'); diff --git a/sample/packedIntegerDotProduct/meta.ts b/sample/packedIntegerDotProduct/meta.ts index d213e688..e20bcac1 100644 --- a/sample/packedIntegerDotProduct/meta.ts +++ b/sample/packedIntegerDotProduct/meta.ts @@ -1,12 +1,7 @@ export default { name: 'Packed Integer Dot Product', description: - 'Compares signed 8-bit dot products using the packed WGSL instruction and an equivalent scalar route, with exact result validation.', + '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: 'reference.ts' }, - { path: 'packed.wgsl' }, - { path: 'scalar.wgsl' }, - ], + sources: [{ path: 'main.ts' }, { path: 'packed.wgsl' }], }; diff --git a/sample/packedIntegerDotProduct/reference.ts b/sample/packedIntegerDotProduct/reference.ts deleted file mode 100644 index 26a70b1a..00000000 --- a/sample/packedIntegerDotProduct/reference.ts +++ /dev/null @@ -1,117 +0,0 @@ -export const packedDotLanguageFeature = 'packed_4x8_integer_dot_product'; - -export type RequestedRoute = 'auto' | 'packed' | 'scalar'; -export type EffectiveRoute = 'packed' | 'scalar'; - -export type RouteSelection = { - requestedRoute: RequestedRoute; - effectiveRoute: EffectiveRoute; - packedLanguageFeatureSupported: boolean; - fallbackReason: string | null; -}; - -function assertI8(value: number) { - if (!Number.isInteger(value) || value < -128 || value > 127) { - throw new RangeError(`expected signed 8-bit integer, received ${value}`); - } -} - -export function pack4xI8(values: readonly number[]): number { - if (values.length !== 4) { - throw new RangeError(`expected 4 values, received ${values.length}`); - } - - let packed = 0; - values.forEach((value, index) => { - assertI8(value); - packed |= (value & 0xff) << (index * 8); - }); - return packed >>> 0; -} - -export function dot4I8(lhs: readonly number[], rhs: readonly number[]): number { - if (lhs.length !== 4 || rhs.length !== 4) { - throw new RangeError('signed dot product requires two 4-component vectors'); - } - - let result = 0; - for (let index = 0; index < 4; ++index) { - assertI8(lhs[index]); - assertI8(rhs[index]); - result += lhs[index] * rhs[index]; - } - return result; -} - -export function selectRoute( - requestedRoute: RequestedRoute, - languageFeatures: ReadonlySet -): RouteSelection { - const packedLanguageFeatureSupported = languageFeatures.has( - packedDotLanguageFeature - ); - const wantsPacked = requestedRoute !== 'scalar'; - const effectiveRoute = - wantsPacked && packedLanguageFeatureSupported ? 'packed' : 'scalar'; - - return { - requestedRoute, - effectiveRoute, - packedLanguageFeatureSupported, - fallbackReason: - requestedRoute === 'packed' && !packedLanguageFeatureSupported - ? `WGSL language feature '${packedDotLanguageFeature}' is unavailable` - : null, - }; -} - -export function validateRun({ - requestedRoute, - effectiveRoute, - packedLanguageFeatureSupported, - expected, - actual, -}: RouteSelection & { - expected: ArrayLike; - actual: ArrayLike; -}) { - if (effectiveRoute === 'packed' && !packedLanguageFeatureSupported) { - throw new Error( - `requested route '${requestedRoute}' claims packed route without language feature support` - ); - } - if (actual.length !== expected.length) { - throw new Error( - `partial output: expected ${expected.length} results, received ${actual.length}` - ); - } - for (let index = 0; index < expected.length; ++index) { - if (actual[index] !== expected[index]) { - throw new Error( - `result ${index} mismatch: expected ${expected[index]}, received ${actual[index]}` - ); - } - } -} - -export function makeInputVectors(count: number) { - const lhsPacked = new Uint32Array(count); - const rhsPacked = new Uint32Array(count); - const expected = new Int32Array(count); - let state = 0x13579bdf; - - const nextI8 = () => { - state = (Math.imul(state, 1664525) + 1013904223) >>> 0; - return ((state >>> 24) & 0xff) - 128; - }; - - for (let index = 0; index < count; ++index) { - const lhs = [nextI8(), nextI8(), nextI8(), nextI8()]; - const rhs = [nextI8(), nextI8(), nextI8(), nextI8()]; - lhsPacked[index] = pack4xI8(lhs); - rhsPacked[index] = pack4xI8(rhs); - expected[index] = dot4I8(lhs, rhs); - } - - return { lhsPacked, rhsPacked, expected }; -} diff --git a/sample/packedIntegerDotProduct/scalar.wgsl b/sample/packedIntegerDotProduct/scalar.wgsl deleted file mode 100644 index 7bf067e4..00000000 --- a/sample/packedIntegerDotProduct/scalar.wgsl +++ /dev/null @@ -1,24 +0,0 @@ -@group(0) @binding(0) var lhs: array; -@group(0) @binding(1) var rhs: array; -@group(0) @binding(2) var output: array; - -fn unpackI8(value: u32, component: u32) -> i32 { - let byte = i32((value >> (component * 8u)) & 0xffu); - return select(byte, byte - 256, byte >= 128); -} - -fn dot4I8Scalar(a: u32, b: u32) -> i32 { - var result = 0; - for (var component = 0u; component < 4u; component += 1u) { - result += unpackI8(a, component) * unpackI8(b, component); - } - return result; -} - -@compute @workgroup_size(64) -fn main(@builtin(global_invocation_id) id: vec3) { - if (id.x >= arrayLength(&output)) { - return; - } - output[id.x] = dot4I8Scalar(lhs[id.x], rhs[id.x]); -} diff --git a/test/packedIntegerDotProduct.test.mjs b/test/packedIntegerDotProduct.test.mjs deleted file mode 100644 index 59376e54..00000000 --- a/test/packedIntegerDotProduct.test.mjs +++ /dev/null @@ -1,93 +0,0 @@ -import assert from 'node:assert/strict'; -import { readFile } from 'node:fs/promises'; -import test from 'node:test'; -import ts from 'typescript'; - -const referenceModuleUrl = new URL( - '../sample/packedIntegerDotProduct/reference.ts', - import.meta.url -); - -async function loadReferenceModule() { - // Node 20 cannot import TypeScript directly. Compile this standalone module - // with the same TypeScript dependency used by the sample build. - const source = await readFile(referenceModuleUrl, 'utf8'); - const { outputText } = ts.transpileModule(source, { - compilerOptions: { - target: ts.ScriptTarget.ES2022, - module: ts.ModuleKind.ESNext, - }, - }); - return import(`data:text/javascript,${encodeURIComponent(outputText)}`); -} - -test('signed packed dot products match scalar arithmetic', async () => { - const { dot4I8, pack4xI8 } = await loadReferenceModule(); - const lhs = [1, -2, 3, -4]; - const rhs = [-5, 6, -7, 8]; - - assert.equal(pack4xI8([1, -2, 127, -128]), 0x807ffe01); - assert.equal(dot4I8(lhs, rhs), -70); -}); - -test('route selection records requested and effective identity', async () => { - const { selectRoute } = await loadReferenceModule(); - const feature = 'packed_4x8_integer_dot_product'; - - assert.deepEqual(selectRoute('auto', new Set([feature])), { - requestedRoute: 'auto', - effectiveRoute: 'packed', - packedLanguageFeatureSupported: true, - fallbackReason: null, - }); - assert.deepEqual(selectRoute('packed', new Set()), { - requestedRoute: 'packed', - effectiveRoute: 'scalar', - packedLanguageFeatureSupported: false, - fallbackReason: `WGSL language feature '${feature}' is unavailable`, - }); - assert.deepEqual(selectRoute('scalar', new Set([feature])), { - requestedRoute: 'scalar', - effectiveRoute: 'scalar', - packedLanguageFeatureSupported: true, - fallbackReason: null, - }); -}); - -test('evidence validation rejects wrong routes and partial output', async () => { - const { validateRun } = await loadReferenceModule(); - - assert.throws( - () => - validateRun({ - requestedRoute: 'packed', - effectiveRoute: 'packed', - packedLanguageFeatureSupported: false, - expected: [11, 22], - actual: [11, 22], - }), - /claims packed route without language feature support/ - ); - assert.throws( - () => - validateRun({ - requestedRoute: 'auto', - effectiveRoute: 'scalar', - packedLanguageFeatureSupported: false, - expected: [11, 22], - actual: [11], - }), - /partial output: expected 2 results, received 1/ - ); - assert.throws( - () => - validateRun({ - requestedRoute: 'auto', - effectiveRoute: 'scalar', - packedLanguageFeatureSupported: false, - expected: [11, 22], - actual: [11, 23], - }), - /result 1 mismatch: expected 22, received 23/ - ); -}); From 0ca514e589c0d0ecd4aabf44a7f3d96cebccb998 Mon Sep 17 00:00:00 2001 From: Kai Ninomiya Date: Wed, 16 Sep 2026 13:59:57 -0700 Subject: [PATCH 4/6] remove gui, tweak input structure --- sample/packedIntegerDotProduct/index.html | 5 -- sample/packedIntegerDotProduct/main.ts | 66 ++++++++++------------ sample/packedIntegerDotProduct/packed.wgsl | 14 +++-- 3 files changed, 39 insertions(+), 46 deletions(-) diff --git a/sample/packedIntegerDotProduct/index.html b/sample/packedIntegerDotProduct/index.html index 615321d8..91fd5ec1 100644 --- a/sample/packedIntegerDotProduct/index.html +++ b/sample/packedIntegerDotProduct/index.html @@ -9,15 +9,10 @@ color-scheme: light dark; } body { - margin: 0; min-height: 100vh; - display: grid; - place-items: center; } pre { - margin: 60px 16px; white-space: pre-wrap; - overflow-wrap: anywhere; } diff --git a/sample/packedIntegerDotProduct/main.ts b/sample/packedIntegerDotProduct/main.ts index 007a71f3..d2ff9ba6 100644 --- a/sample/packedIntegerDotProduct/main.ts +++ b/sample/packedIntegerDotProduct/main.ts @@ -1,7 +1,15 @@ -import { GUI } from 'dat.gui'; import packedWGSL from './packed.wgsl'; import { quitIfWebGPUNotAvailableOrMissingFeatures } from '../util'; +const kSampleCases = [ + { lhs: [1, -2, 3, -4], rhs: [-5, 6, -7, 8] }, + { lhs: [-128, -128, -128, -128], rhs: [-128, -128, -128, -128] }, + { lhs: [127, 127, 127, 127], rhs: [127, 127, 127, 127] }, + { lhs: [1, 2, 3, 4], rhs: [4, 3, 2, 1] }, +] as const; + +const kWorkgroupSize = 64; // Same as in packed.wgsl + const result = document.querySelector('#result') as HTMLElement; if ( !navigator.gpu?.wgslLanguageFeatures.has('packed_4x8_integer_dot_product') @@ -15,20 +23,9 @@ if ( const device = await adapter?.requestDevice(); quitIfWebGPUNotAvailableOrMissingFeatures(adapter, device); - const lhs = [ - [1, -2, 3, -4], - [-128, -128, -128, -128], - [127, 127, 127, 127], - [1, 2, 3, 4], - ]; - const rhs = [ - [-5, 6, -7, 8], - [-128, -128, -128, -128], - [127, 127, 127, 127], - [4, 3, 2, 1], - ]; - - function createInputBuffer(vectors: number[][]) { + function createInputBuffer( + vectors: ReadonlyArray + ) { // Pack four signed 8-bit components into each u32, low byte first. // Masking preserves the two's-complement representation of negative values. const packed = new Uint32Array( @@ -50,7 +47,7 @@ if ( return buffer; } - const outputSize = lhs.length * Int32Array.BYTES_PER_ELEMENT; + const outputSize = kSampleCases.length * Int32Array.BYTES_PER_ELEMENT; const outputBuffer = device.createBuffer({ size: outputSize, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC, @@ -63,12 +60,14 @@ if ( layout: 'auto', compute: { module: device.createShaderModule({ code: packedWGSL }) }, }); + const inputBuffer = createInputBuffer( + kSampleCases.flatMap((c) => [c.lhs, c.rhs]) + ); const bindGroup = device.createBindGroup({ layout: pipeline.getBindGroupLayout(0), entries: [ - { binding: 0, resource: { buffer: createInputBuffer(lhs) } }, - { binding: 1, resource: { buffer: createInputBuffer(rhs) } }, - { binding: 2, resource: { buffer: outputBuffer } }, + { binding: 0, resource: { buffer: inputBuffer } }, + { binding: 1, resource: { buffer: outputBuffer } }, ], }); @@ -76,7 +75,7 @@ if ( const pass = encoder.beginComputePass(); pass.setPipeline(pipeline); pass.setBindGroup(0, bindGroup); - pass.dispatchWorkgroups(Math.ceil(lhs.length / 64)); + pass.dispatchWorkgroups(Math.ceil(kSampleCases.length / kWorkgroupSize)); pass.end(); encoder.copyBufferToBuffer(outputBuffer, 0, readbackBuffer, 0, outputSize); device.queue.submit([encoder.finish()]); @@ -85,21 +84,16 @@ if ( const results = new Int32Array(readbackBuffer.getMappedRange()).slice(); readbackBuffer.unmap(); - const settings = { example: 0 }; - function showResult() { - const i = settings.example; - result.textContent = `a = [${lhs[i].join(', ')}] -b = [${rhs[i].join(', ')}] -dot4I8Packed(a, b) = ${results[i]}`; + for (const [i, sample] of kSampleCases.entries()) { + const expected = + sample.lhs[0] * sample.rhs[0] + + sample.lhs[1] * sample.rhs[1] + + sample.lhs[2] * sample.rhs[2] + + sample.lhs[3] * sample.rhs[3]; + const lhs = sample.lhs.map((x) => x.toString().padStart(4)).join(', '); + const rhs = sample.rhs.map((x) => x.toString().padStart(4)).join(', '); + const out = results[i].toString().padStart(6); + const exp = expected.toString().padStart(6); + result.textContent += `\ndot4I8Packed of [${lhs}] by [${rhs}] gave ${out} (expecting ${exp})`; } - const gui = new GUI(); - gui - .add(settings, 'example', { - 'Mixed signs': 0, - 'Minimum signed bytes': 1, - 'Maximum signed bytes': 2, - 'Positive components': 3, - }) - .onChange(showResult); - showResult(); } diff --git a/sample/packedIntegerDotProduct/packed.wgsl b/sample/packedIntegerDotProduct/packed.wgsl index 46b68fde..8a370afd 100644 --- a/sample/packedIntegerDotProduct/packed.wgsl +++ b/sample/packedIntegerDotProduct/packed.wgsl @@ -1,13 +1,17 @@ requires packed_4x8_integer_dot_product; -@group(0) @binding(0) var lhs: array; -@group(0) @binding(1) var rhs: array; -@group(0) @binding(2) var output: array; +struct Case { lhs: u32, rhs: u32 } -@compute @workgroup_size(64) +@group(0) @binding(0) var input: array; +@group(0) @binding(1) var output: array; + +const kWorkgroupSize: u32 = 64; // Same as in main.ts + +@compute @workgroup_size(kWorkgroupSize) fn main(@builtin(global_invocation_id) id: vec3) { if (id.x >= arrayLength(&output)) { return; } - output[id.x] = dot4I8Packed(lhs[id.x], rhs[id.x]); + + output[id.x] = dot4I8Packed(input[id.x].lhs, input[id.x].rhs); } From 0495650eb01c2213ec50a4562452dba49abbb5fb Mon Sep 17 00:00:00 2001 From: Kai Ninomiya Date: Wed, 16 Sep 2026 14:14:16 -0700 Subject: [PATCH 5/6] tweaks --- sample/packedIntegerDotProduct/index.html | 2 +- sample/packedIntegerDotProduct/main.ts | 43 +++++++++++++--------- sample/packedIntegerDotProduct/packed.wgsl | 2 +- 3 files changed, 27 insertions(+), 20 deletions(-) diff --git a/sample/packedIntegerDotProduct/index.html b/sample/packedIntegerDotProduct/index.html index 91fd5ec1..a44b87f5 100644 --- a/sample/packedIntegerDotProduct/index.html +++ b/sample/packedIntegerDotProduct/index.html @@ -9,7 +9,7 @@ color-scheme: light dark; } body { - min-height: 100vh; + min-height: 20em; } pre { white-space: pre-wrap; diff --git a/sample/packedIntegerDotProduct/main.ts b/sample/packedIntegerDotProduct/main.ts index d2ff9ba6..d3f73977 100644 --- a/sample/packedIntegerDotProduct/main.ts +++ b/sample/packedIntegerDotProduct/main.ts @@ -8,7 +8,17 @@ const kSampleCases = [ { lhs: [1, 2, 3, 4], rhs: [4, 3, 2, 1] }, ] as const; -const kWorkgroupSize = 64; // Same as in packed.wgsl +const kWorkgroupSize = 64; // Must match packed.wgsl + +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. + return ( + (x & 0xff) | ((y & 0xff) << 8) | ((z & 0xff) << 16) | ((w & 0xff) << 24) + ) >>> 0; +} const result = document.querySelector('#result') as HTMLElement; if ( @@ -23,20 +33,8 @@ if ( const device = await adapter?.requestDevice(); quitIfWebGPUNotAvailableOrMissingFeatures(adapter, device); - function createInputBuffer( - vectors: ReadonlyArray - ) { - // Pack four signed 8-bit components into each u32, low byte first. - // Masking preserves the two's-complement representation of negative values. - const packed = new Uint32Array( - vectors.map( - ([x, y, z, w]) => - (x & 0xff) | - ((y & 0xff) << 8) | - ((z & 0xff) << 16) | - ((w & 0xff) << 24) - ) - ); + function createInputBuffer(vectors: vec4i[]) { + const packed = new Uint32Array(vectors.map(pack4xI8)); const buffer = device.createBuffer({ size: packed.byteLength, usage: GPUBufferUsage.STORAGE, @@ -85,15 +83,24 @@ if ( readbackBuffer.unmap(); for (const [i, sample] of kSampleCases.entries()) { + // Result should be the same in JS, show that for comparison. const expected = sample.lhs[0] * sample.rhs[0] + sample.lhs[1] * sample.rhs[1] + sample.lhs[2] * sample.rhs[2] + sample.lhs[3] * sample.rhs[3]; - const lhs = sample.lhs.map((x) => x.toString().padStart(4)).join(', '); - const rhs = sample.rhs.map((x) => x.toString().padStart(4)).join(', '); + + const lhs = `[${sample.lhs + .map((x) => x.toString().padStart(4)) + .join(', ')}] (0x${pack4xI8(sample.lhs).toString(16).padStart(8, '0')})`; + const rhs = `[${sample.rhs + .map((x) => x.toString().padStart(4)) + .join(', ')}] (0x${pack4xI8(sample.rhs).toString(16).padStart(8, '0')})`; const out = results[i].toString().padStart(6); const exp = expected.toString().padStart(6); - result.textContent += `\ndot4I8Packed of [${lhs}] by [${rhs}] gave ${out} (expecting ${exp})`; + result.textContent += ` + +WGSL dot4I8Packed of ${lhs} + by ${rhs} gave ${out} (JS gave ${exp})`; } } diff --git a/sample/packedIntegerDotProduct/packed.wgsl b/sample/packedIntegerDotProduct/packed.wgsl index 8a370afd..253e27ba 100644 --- a/sample/packedIntegerDotProduct/packed.wgsl +++ b/sample/packedIntegerDotProduct/packed.wgsl @@ -5,7 +5,7 @@ struct Case { lhs: u32, rhs: u32 } @group(0) @binding(0) var input: array; @group(0) @binding(1) var output: array; -const kWorkgroupSize: u32 = 64; // Same as in main.ts +const kWorkgroupSize: u32 = 64; // Must match main.ts @compute @workgroup_size(kWorkgroupSize) fn main(@builtin(global_invocation_id) id: vec3) { From 57577d6089193af7256091ef5c9f42963a7bf805 Mon Sep 17 00:00:00 2001 From: Kai Ninomiya Date: Wed, 16 Sep 2026 14:38:57 -0700 Subject: [PATCH 6/6] add a gui with the actual lhs/rhs values to make it interactive --- sample/packedIntegerDotProduct/index.html | 7 +- sample/packedIntegerDotProduct/main.ts | 129 ++++++++++++--------- sample/packedIntegerDotProduct/packed.wgsl | 18 +-- 3 files changed, 85 insertions(+), 69 deletions(-) diff --git a/sample/packedIntegerDotProduct/index.html b/sample/packedIntegerDotProduct/index.html index a44b87f5..0467ed5a 100644 --- a/sample/packedIntegerDotProduct/index.html +++ b/sample/packedIntegerDotProduct/index.html @@ -9,16 +9,19 @@ color-scheme: light dark; } body { - min-height: 20em; + width: calc(100vw - 280px); + height: 400px; + overflow-y: hidden; } pre { white-space: pre-wrap; + height: 370px; } - 
Computing packed integer dot products…
+ 
Computing packed integer dot products…
diff --git a/sample/packedIntegerDotProduct/main.ts b/sample/packedIntegerDotProduct/main.ts index d3f73977..ef753e73 100644 --- a/sample/packedIntegerDotProduct/main.ts +++ b/sample/packedIntegerDotProduct/main.ts @@ -1,26 +1,20 @@ +import { GUI } from 'dat.gui'; import packedWGSL from './packed.wgsl'; import { quitIfWebGPUNotAvailableOrMissingFeatures } from '../util'; -const kSampleCases = [ - { lhs: [1, -2, 3, -4], rhs: [-5, 6, -7, 8] }, - { lhs: [-128, -128, -128, -128], rhs: [-128, -128, -128, -128] }, - { lhs: [127, 127, 127, 127], rhs: [127, 127, 127, 127] }, - { lhs: [1, 2, 3, 4], rhs: [4, 3, 2, 1] }, -] as const; - -const kWorkgroupSize = 64; // Must match packed.wgsl - 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. - return ( - (x & 0xff) | ((y & 0xff) << 8) | ((z & 0xff) << 16) | ((w & 0xff) << 24) - ) >>> 0; + /*prettier-ignore*/ + return ((x & 0xff) | + ((y & 0xff) << 8) | + ((z & 0xff) << 16) | + ((w & 0xff) << 24)) >>> 0; } -const result = document.querySelector('#result') as HTMLElement; +const outputElement = document.querySelector('#output') as HTMLElement; if ( !navigator.gpu?.wgslLanguageFeatures.has('packed_4x8_integer_dot_product') ) { @@ -33,34 +27,26 @@ if ( const device = await adapter?.requestDevice(); quitIfWebGPUNotAvailableOrMissingFeatures(adapter, device); - function createInputBuffer(vectors: vec4i[]) { - const packed = new Uint32Array(vectors.map(pack4xI8)); - const buffer = device.createBuffer({ - size: packed.byteLength, - usage: GPUBufferUsage.STORAGE, - mappedAtCreation: true, - }); - new Uint32Array(buffer.getMappedRange()).set(packed); - buffer.unmap(); - return buffer; - } + const kInputSize = 2 * Uint32Array.BYTES_PER_ELEMENT; + const inputBuffer = device.createBuffer({ + size: kInputSize, + usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.STORAGE, + }); - const outputSize = kSampleCases.length * Int32Array.BYTES_PER_ELEMENT; + const kOutputSize = Int32Array.BYTES_PER_ELEMENT; const outputBuffer = device.createBuffer({ - size: outputSize, + size: kOutputSize, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_SRC, }); const readbackBuffer = device.createBuffer({ - size: outputSize, + size: kOutputSize, usage: GPUBufferUsage.COPY_DST | GPUBufferUsage.MAP_READ, }); + const pipeline = await device.createComputePipelineAsync({ layout: 'auto', compute: { module: device.createShaderModule({ code: packedWGSL }) }, }); - const inputBuffer = createInputBuffer( - kSampleCases.flatMap((c) => [c.lhs, c.rhs]) - ); const bindGroup = device.createBindGroup({ layout: pipeline.getBindGroupLayout(0), entries: [ @@ -69,38 +55,71 @@ if ( ], }); - const encoder = device.createCommandEncoder(); - const pass = encoder.beginComputePass(); - pass.setPipeline(pipeline); - pass.setBindGroup(0, bindGroup); - pass.dispatchWorkgroups(Math.ceil(kSampleCases.length / kWorkgroupSize)); - pass.end(); - encoder.copyBufferToBuffer(outputBuffer, 0, readbackBuffer, 0, outputSize); - device.queue.submit([encoder.finish()]); + async function updateResult() { + // If an update is still in progress just wait until it's done. + if (readbackBuffer.mapState !== 'unmapped') { + setTimeout(updateResult, 0); + return; + } - await readbackBuffer.mapAsync(GPUMapMode.READ); - const results = new Int32Array(readbackBuffer.getMappedRange()).slice(); - readbackBuffer.unmap(); + 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]; - for (const [i, sample] of kSampleCases.entries()) { // Result should be the same in JS, show that for comparison. const expected = - sample.lhs[0] * sample.rhs[0] + - sample.lhs[1] * sample.rhs[1] + - sample.lhs[2] * sample.rhs[2] + - sample.lhs[3] * sample.rhs[3]; + lhs[0] * rhs[0] + lhs[1] * rhs[1] + lhs[2] * rhs[2] + lhs[3] * rhs[3]; - const lhs = `[${sample.lhs + const lhsStr = `[${lhs .map((x) => x.toString().padStart(4)) - .join(', ')}] (0x${pack4xI8(sample.lhs).toString(16).padStart(8, '0')})`; - const rhs = `[${sample.rhs + .join(', ')}] (0x${pack4xI8(lhs).toString(16).padStart(8, '0')})`; + const rhsStr = `[${rhs .map((x) => x.toString().padStart(4)) - .join(', ')}] (0x${pack4xI8(sample.rhs).toString(16).padStart(8, '0')})`; - const out = results[i].toString().padStart(6); - const exp = expected.toString().padStart(6); - result.textContent += ` + .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 ${lhs} - by ${rhs} gave ${out} (JS gave ${exp})`; +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(); } diff --git a/sample/packedIntegerDotProduct/packed.wgsl b/sample/packedIntegerDotProduct/packed.wgsl index 253e27ba..3ae0001d 100644 --- a/sample/packedIntegerDotProduct/packed.wgsl +++ b/sample/packedIntegerDotProduct/packed.wgsl @@ -1,17 +1,11 @@ requires packed_4x8_integer_dot_product; -struct Case { lhs: u32, rhs: u32 } +struct Input { lhs: u32, rhs: u32 } -@group(0) @binding(0) var input: array; -@group(0) @binding(1) var output: array; +@group(0) @binding(0) var input: Input; +@group(0) @binding(1) var output: i32; -const kWorkgroupSize: u32 = 64; // Must match main.ts - -@compute @workgroup_size(kWorkgroupSize) -fn main(@builtin(global_invocation_id) id: vec3) { - if (id.x >= arrayLength(&output)) { - return; - } - - output[id.x] = dot4I8Packed(input[id.x].lhs, input[id.x].rhs); +@compute @workgroup_size(1) +fn main() { + output = dot4I8Packed(input.lhs, input.rhs); }