Skip to content

Instantly share code, notes, and snippets.

@bellbind
Last active May 9, 2022
Embed
What would you like to do?
[WebGPU] Game of Life as Compute+Render example for Chrome-103
<!doctype html>
<html>
<head>
<!-- IMPORTANT: The current Chrome requires some origin-trial token in <meta>.
To register origins at the last "WebGPU REGISTER" in https://developer.chrome.com/origintrials/
This token is for a Web Origin "https://gist.githack.com" (maybe expired at Mar 31, 2022)
It can register localhost origin as "http://localhost:8000"
-->
<meta http-equiv="origin-trial"
content="Akv07qcAop5MFaZYxJtHHjUuM8eV3GpbHkTeuhZo/4wsNjYnQ7GSGJyo7hRVZvpvyjYwilbJ8KbFVchI4O1DpA0AAABQeyJvcmlnaW4iOiJodHRwczovL2dpc3QuZ2l0aGFjay5jb206NDQzIiwiZmVhdHVyZSI6IldlYkdQVSIsImV4cGlyeSI6MTY1MjgzMTk5OX0=" />
<meta http-equiv="origin-trial"
content="AkIL+/THBoi1QEsWbX5SOuMpL6+KGAXKrZE5Bz6yHTuijzvKz2MznuLqE+MH4YSqRi/v1fDK/6JyFzgibTTeNAsAAABJeyJvcmlnaW4iOiJodHRwOi8vbG9jYWxob3N0OjgwMDAiLCJmZWF0dXJlIjoiV2ViR1BVIiwiZXhwaXJ5IjoxNjUyODMxOTk5fQ==" />
<script src="./main.js" type="module"></script>
<style>@media(prefers-color-scheme: dark){:root {color-scheme: dark;}}</style>
<link rel="icon" href="data:image/x-icon;," />
</head>
<body>
<h1>(Notice: The origin-trial token in this page will be expired at May 15, 2022)</h1>
<canvas style="width: 80vmin; height: 80vmin; border: solid;" id="canvas"></canvas>
</body>
</html>
// Lifegame as a simple Compute+Render example for WebGPU API for Chrome-103: https://www.w3.org/TR/webgpu/
const adapter = await navigator.gpu.requestAdapter();
const device = await adapter.requestDevice();
//[compute setup]
// buffers of cell tables
const width = 256, height = 256;
const cells = new Uint32Array(width * height);
for (let i = 0; i < 10000;) {
const n = Math.trunc(Math.random() * width * height);
if (cells[n] === 1) continue;
cells[n] = 1;
i++;
}
const buffer0 = device.createBuffer({size: cells.byteLength, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.VERTEX, mappedAtCreation: true});
new Uint32Array(buffer0.getMappedRange()).set(cells);
buffer0.unmap();
const buffer1 = device.createBuffer({size: cells.byteLength, usage: GPUBufferUsage.STORAGE | GPUBufferUsage.VERTEX});
// compute shader of step
const workgroupSize = adapter.limits.maxComputeInvocationsPerWorkgroup ?
2 ** Math.trunc(Math.log2(Math.sqrt(adapter.limits.maxComputeInvocationsPerWorkgroup))) : 16;
const computeWgsl = `
// 'override' constants after implemented
let w = ${width};
let h = ${height};
let ws = ${workgroupSize};
struct Cells {
data: array<u32>,
};
@binding(0) @group(0) var<storage, read> current: Cells;
@binding(1) @group(0) var<storage, write> next: Cells;
fn index(x: i32, y: i32) -> u32 {
return u32(((y + h) % h) * w + ((x + w) % w));
}
fn cell(x: i32, y: i32) -> u32 {
return current.data[index(x, y)];
}
fn neighbors(x: i32, y: i32) -> u32 {
return cell(x - 1, y - 1) + cell(x, y - 1) + cell(x + 1, y - 1) +
cell(x - 1, y) + cell(x + 1, y) +
cell(x - 1, y + 1) + cell(x, y + 1) + cell(x + 1, y + 1);
}
@stage(compute) @workgroup_size(ws, ws) fn step_next(@builtin(global_invocation_id) giid: vec3<u32>) {
let x = i32(giid.x);
let y = i32(giid.y);
let n = neighbors(x, y);
next.data[index(x, y)] = select(u32(n == 3u), u32(n == 2u || n == 3u), cell(x, y) == 1u);
}
`;
const computeShader = device.createShaderModule({code: computeWgsl});
// compute pipeline
const computePipeline = device.createComputePipeline({
compute: {module: computeShader, entryPoint: "step_next", constants: {w: width, h: height, ws: workgroupSize}},
});
const bindGroupLayout = computePipeline.getBindGroupLayout(0);
const bindGroup0 = device.createBindGroup({
layout: bindGroupLayout,
entries: [
{binding: 0, resource: {buffer: buffer0}},
{binding: 1, resource: {buffer: buffer1}},
]
});
const bindGroup1 = device.createBindGroup({
layout: bindGroupLayout,
entries: [
{binding: 0, resource: {buffer: buffer1}},
{binding: 1, resource: {buffer: buffer0}},
]
});
//[render setup]
// cell vertex buffer
const corners = new Uint32Array([0, 0, 0, 1, 1, 0, 1, 1]); // 4-offset of (x, y)
const cornersBuffer = device.createBuffer({size: corners.byteLength, usage: GPUBufferUsage.VERTEX, mappedAtCreation: true});
new Uint32Array(cornersBuffer.getMappedRange()).set(corners);
cornersBuffer.unmap();
//const cornersBuffer = device.createBuffer({size: corners.byteLength, usage: GPUBufferUsage.VERTEX | GPUBufferUsage.COPY_DST});
//device.queue.writeBuffer(cornersBuffer, 0, corners.buffer, corners.byteOffset, corners.byteLength);
const cornersStride = {arrayStride: 2 * corners.BYTES_PER_ELEMENT, stepMode: "vertex", attributes: [{shaderLocation: 1, offset: 0, format: "uint32x2"}]};
// stride for buffer0/buffer1 as instances
const cellsStride = {arrayStride: cells.BYTES_PER_ELEMENT, stepMode: "instance", attributes: [{shaderLocation: 0, offset: 0, format: "uint32"}]};
// cell shader
const vertexWgsl = `
struct Out {
@builtin(position) pos: vec4<f32>,
@location(0) cell: f32,
};
@stage(vertex) fn main(@builtin(instance_index) i: u32, @location(0) cell: u32, @location(1) v: vec2<u32>) -> Out {
let w = ${width}u;
let h = ${height}u;
let x = (f32(i % w + v.x) / f32(w) - 0.5) * 2.0;
let y = (f32((i - (i % w)) / w + v.y) / f32(h) - 0.5) * 2.0;
return Out(vec4<f32>(x, y, 0.0, 1.0), f32(cell));
}
`;
const vertexShader = device.createShaderModule({code: vertexWgsl});
const fragmentWgsl = `
@stage(fragment) fn main(@location(0) cell: f32) -> @location(0) vec4<f32> {
return vec4<f32>(f32(cell), f32(cell), f32(cell), 1.0);
}
`;
const fragmentShader = device.createShaderModule({code: fragmentWgsl});
// gpu config for canvas
const canvas = document.getElementById("canvas");
const gpu = canvas.getContext("webgpu");
const format = gpu.getPreferredFormat(adapter);
const compositingAlphaMode = "premultiplied";
gpu.configure({device, format, compositingAlphaMode, size: [canvas.width, canvas.height]});
// render pipeline
const renderPipeline = device.createRenderPipeline({
primitive: {topology: "triangle-strip"},
vertex: {module: vertexShader, entryPoint: "main", buffers: [cellsStride, cornersStride]},
fragment: {module: fragmentShader, entryPoint: "main", targets: [{format}]},
});
//[command part]
// render
const render = (t, compute = true, render = true) => {
const view = gpu.getCurrentTexture().createView();
const renderPass = {colorAttachments: [{view, loadOp: "clear", storeOp: "store"}]};
const commandEncoder = device.createCommandEncoder();
if (compute) {
const passEncoder = commandEncoder.beginComputePass();
passEncoder.setPipeline(computePipeline);
passEncoder.setBindGroup(0, t % 2 === 0 ? bindGroup0 : bindGroup1);
passEncoder.dispatchWorkgroups(width / workgroupSize, height / workgroupSize);
passEncoder.end();
}
if (render) {
const passEncoder = commandEncoder.beginRenderPass(renderPass);
passEncoder.setPipeline(renderPipeline);
passEncoder.setVertexBuffer(0, t % 2 === 0 ? buffer1 : buffer0);
passEncoder.setVertexBuffer(1, cornersBuffer);
passEncoder.draw(4, width * height);
passEncoder.end();
}
device.queue.submit([commandEncoder.finish()]);
};
(function loop(t) {
render(t);
requestAnimationFrame(() => loop(t ^ 1));
})(0);
@bellbind
Copy link
Author

bellbind commented Mar 4, 2022

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment