bellbind/index.html

## index.html
<!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>

## main.js
// Game of Life as a simple Compute+Render example for WebGPU API: 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 blockAttr = navigator.userAgent?.match(/Chrome\/98/) ? "[[block]]" : "";
const workgroupSize = adapter.limits.maxComputeInvocationsPerWorkgroup ?
      2 ** Math.trunc(Math.log2(Math.sqrt(adapter.limits.maxComputeInvocationsPerWorkgroup))) : 16;
const computeWgsl = `
${blockAttr} 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 {
  let w = ${width};
  let h = ${height};
  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(${workgroupSize}, ${workgroupSize})]] 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"},
});
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 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);
gpu.configure({device, format, 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, loadValue: {r: 0, g: 0, b: 0, a: 0}, storeOp: "store"}]};
  const commandEncoder = device.createCommandEncoder();
  if (compute) {
    const passEncoder = commandEncoder.beginComputePass();
    passEncoder.setPipeline(computePipeline);
    passEncoder.setBindGroup(0, t % 2 === 0 ? bindGroup0 : bindGroup1);
    passEncoder.dispatch(width / workgroupSize, height / workgroupSize);
    passEncoder.endPass();
  }
  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.endPass();
  }
  device.queue.submit([commandEncoder.finish()]);
};
(function loop(t) {
  render(t);
  requestAnimationFrame(() => loop(t ^ 1));
})(0);
	<!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>
	// Game of Life as a simple Compute+Render example for WebGPU API: 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 blockAttr = navigator.userAgent?.match(/Chrome\/98/) ? "[[block]]" : "";
	const workgroupSize = adapter.limits.maxComputeInvocationsPerWorkgroup ?
	2 ** Math.trunc(Math.log2(Math.sqrt(adapter.limits.maxComputeInvocationsPerWorkgroup))) : 16;
	const computeWgsl = `
	${blockAttr} 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 {
	let w = ${width};
	let h = ${height};
	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(${workgroupSize}, ${workgroupSize})]] 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"},
	});
	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 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);
	gpu.configure({device, format, 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, loadValue: {r: 0, g: 0, b: 0, a: 0}, storeOp: "store"}]};
	const commandEncoder = device.createCommandEncoder();
	if (compute) {
	const passEncoder = commandEncoder.beginComputePass();
	passEncoder.setPipeline(computePipeline);
	passEncoder.setBindGroup(0, t % 2 === 0 ? bindGroup0 : bindGroup1);
	passEncoder.dispatch(width / workgroupSize, height / workgroupSize);
	passEncoder.endPass();
	}
	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.endPass();
	}
	device.queue.submit([commandEncoder.finish()]);
	};
	(function loop(t) {
	render(t);
	requestAnimationFrame(() => loop(t ^ 1));
	})(0);