diff --git a/sample/packedIntegerDotProduct/index.html b/sample/packedIntegerDotProduct/index.html new file mode 100644 index 00000000..0467ed5a --- /dev/null +++ b/sample/packedIntegerDotProduct/index.html @@ -0,0 +1,27 @@ + + + + + + webgpu-samples: packedIntegerDotProduct + + + + + + 
Computing packed integer dot products…
+ + diff --git a/sample/packedIntegerDotProduct/main.ts b/sample/packedIntegerDotProduct/main.ts new file mode 100644 index 00000000..ef753e73 --- /dev/null +++ b/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(); +} diff --git a/sample/packedIntegerDotProduct/meta.ts b/sample/packedIntegerDotProduct/meta.ts new file mode 100644 index 00000000..e20bcac1 --- /dev/null +++ b/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' }], +}; diff --git a/sample/packedIntegerDotProduct/packed.wgsl b/sample/packedIntegerDotProduct/packed.wgsl new file mode 100644 index 00000000..3ae0001d --- /dev/null +++ b/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 input: Input; +@group(0) @binding(1) var output: i32; + +@compute @workgroup_size(1) +fn main() { + output = dot4I8Packed(input.lhs, input.rhs); +} 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, }, },