Skip to content

Instantly share code, notes, and snippets.

@shogonir
Created January 4, 2024 17:55
Show Gist options
  • Save shogonir/c5ee2383705f942c4f64504597b18a8f to your computer and use it in GitHub Desktop.
Save shogonir/c5ee2383705f942c4f64504597b18a8f to your computer and use it in GitHub Desktop.
<!DOCTYPE html>
<html>
<head>
<meta charset="UTF-8">
<title>webgpu-test</title>
</head>
<body>
<canvas id="canvas"></canvas>
<script>
const GRID_SIZE = 32;
const WORKGROUP_SIZE = 8;
const UPDATE_INTERVAL = 200;
let step = 0;
const main = (
canvas,
webgpu,
adapter,
device
) => {
const canvasFormat = navigator.gpu.getPreferredCanvasFormat();
webgpu.configure({
device,
format: canvasFormat,
});
// geometry
const positive = 0.75;
const negative = -0.75;
const vertices = new Float32Array([
// bottom right
negative, negative,
positive, negative,
positive, positive,
// top left
negative, negative,
positive, positive,
negative, positive
]);
// 長さや用途は変更できない。長さが同じで内容を変えるだけの場合は書き換え可能。
const vertexBuffer = device.createBuffer({
label: 'Cell vertices',
size: vertices.byteLength,
usage: GPUBufferUsage.VERTEX | GPUBufferUsage.COPY_DST,
});
const vertexBufferOffset = 0;
device.queue.writeBuffer(vertexBuffer, vertexBufferOffset, vertices);
const vertexBufferLayout = {
arrayStride: 2 * Float32Array.BYTES_PER_ELEMENT, // 単位はバイト。x,yの2要素 * 4バイト
attributes: [
{
format: 'float32x2',
offset: 0,
shaderLocation: 0, // shader の @location(0) に対応
},
],
};
const simulationShaderModule = device.createShaderModule({
label: 'Game of Life simulation shader',
code: `
@group(0) @binding(0) var<uniform> grid: vec2f;
@group(0) @binding(1) var<storage> cellStateIn: array<u32>;
@group(0) @binding(2) var<storage, read_write> cellStateOut: array<u32>;
fn cellIndex(cell: vec2u) -> u32 {
return (cell.y % u32(grid.y)) * u32(grid.x) +
(cell.x % u32(grid.x));
}
fn cellActive(x: u32, y: u32) -> u32 {
return cellStateIn[cellIndex(vec2(x, y))];
}
@compute
@workgroup_size(${WORKGROUP_SIZE}, ${WORKGROUP_SIZE})
fn computeMain(
@builtin(global_invocation_id) cell: vec3u
) {
let activeNeighbors =
cellActive(cell.x + 1, cell.y + 1) +
cellActive(cell.x + 1, cell.y + 0) +
cellActive(cell.x + 1, cell.y - 1) +
cellActive(cell.x + 0, cell.y - 1) +
cellActive(cell.x - 1, cell.y - 1) +
cellActive(cell.x - 1, cell.y + 0) +
cellActive(cell.x - 1, cell.y + 1) +
cellActive(cell.x + 0, cell.y + 1);
let i = cellIndex(cell.xy);
// Conway's game of life rules:
switch activeNeighbors {
case 2: {
cellStateOut[i] = cellStateIn[i];
}
case 3: {
cellStateOut[i] = 1;
}
default: {
cellStateOut[i] = 0;
}
}
}
`,
});
const cellShaderModule = device.createShaderModule({
label: 'Cell shader',
code:
`
struct VertexInput {
@location(0) pos: vec2f,
@builtin(instance_index) instance: u32
};
struct VertexOutput {
@builtin(position) pos: vec4f,
@location(0) cell: vec2f,
};
@group(0) @binding(0) var<uniform> grid: vec2f;
@group(0) @binding(1) var<storage> cellState: array<u32>;
@vertex
fn vertexMain(input: VertexInput) -> VertexOutput {
let i = f32(input.instance);
let cell = vec2f(i % grid.x, floor(i / grid.x));
let state = f32(cellState[input.instance]);
let cellOffset = cell / grid * 2;
let gridPos = (input.pos * state + 1) / grid - 1 + cellOffset;
var output: VertexOutput;
output.pos = vec4f(gridPos, 0, 1);
output.cell = cell;
return output;
}
@fragment
fn fragmentMain(input: VertexOutput) -> @location(0) vec4f {
let c = input.cell / grid;
return vec4f(c, 1 - c.x, 1);
}
`,
});
// uniform
const uniformArray = new Float32Array([GRID_SIZE, GRID_SIZE]);
const uniformBuffer = device.createBuffer({
label: 'Grid Uniforms',
size: uniformArray.byteLength,
usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST,
});
const uniformBufferOffset = 0;
device.queue.writeBuffer(uniformBuffer, uniformBufferOffset, uniformArray);
// storage
const cellStateArray = new Uint32Array(GRID_SIZE * GRID_SIZE);
const cellStateStorages = [
device.createBuffer({
label: 'Cell State A',
size: cellStateArray.byteLength,
usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST,
}),
device.createBuffer({
label: 'Cell State B',
size: cellStateArray.byteLength,
usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST,
}),
];
const storageBufferOffset = 0;
for (let i = 0; i < cellStateArray.length; i++) {
cellStateArray[i] = Math.random() > 0.6 ? 1 : 0;
}
device.queue.writeBuffer(cellStateStorages[0], storageBufferOffset, cellStateArray);
const bindGroupLayout = device.createBindGroupLayout({
label: 'Cell Bind Group Layout',
entries: [
{
binding: 0,
visibility: GPUShaderStage.VERTEX | GPUShaderStage.COMPUTE | GPUShaderStage.FRAGMENT,
buffer: {
// type: 'uniform',
},
},
{
binding: 1,
visibility: GPUShaderStage.VERTEX | GPUShaderStage.COMPUTE,
buffer: {
type: 'read-only-storage',
},
},
{
binding: 2,
visibility: GPUShaderStage.COMPUTE | GPUShaderStage.FRAGMENT,
buffer: {
type: 'storage', // read write storage
},
},
],
});
const pipelineLayout = device.createPipelineLayout({
label: 'Cell Pipeline Layout',
bindGroupLayouts: [
bindGroupLayout,
],
});
const simulationPipeline = device.createComputePipeline({
label: 'Simulation pipeline',
layout: pipelineLayout,
compute: {
module: simulationShaderModule,
entryPoint: 'computeMain',
},
});
const cellPipeline = device.createRenderPipeline({
label: 'Cell pipeline',
layout: pipelineLayout,
vertex: {
module: cellShaderModule,
entryPoint: 'vertexMain',
buffers: [
vertexBufferLayout
],
},
fragment: {
module: cellShaderModule,
entryPoint: 'fragmentMain',
targets: [
{
format: canvasFormat,
}
]
}
});
const bindGroups = [
device.createBindGroup({
label: 'Cell render bind group A',
layout: bindGroupLayout,
entries: [
// uniform
{
binding: 0,
resource: {
buffer: uniformBuffer,
},
},
// storage
{
binding: 1,
resource: {
buffer: cellStateStorages[0],
},
},
{
binding: 2,
resource: {
buffer: cellStateStorages[1],
},
},
],
}),
device.createBindGroup({
label: 'Cell render bind group B',
layout: bindGroupLayout,
entries: [
// uniform
{
binding: 0,
resource: {
buffer: uniformBuffer,
},
},
// storage
{
binding: 1,
resource: {
buffer: cellStateStorages[1],
},
},
{
binding: 2,
resource: {
buffer: cellStateStorages[0],
},
},
],
})
];
const updateGrid = () => {
const encoder = device.createCommandEncoder();
// compute shader
const computePass = encoder.beginComputePass();
computePass.setPipeline(simulationPipeline);
computePass.setBindGroup(0, bindGroups[step % 2]);
const workgroupCount = Math.ceil(GRID_SIZE / WORKGROUP_SIZE);
computePass.dispatchWorkgroups(workgroupCount, workgroupCount);
computePass.end();
step++;
// vertex shader, fragment shader
const pass = encoder.beginRenderPass({
colorAttachments: [
{
view: webgpu.getCurrentTexture().createView(),
loadOp: 'clear',
clearValue: {
r: 0,
g: 0,
b: 0.4,
a: 1,
},
storeOp: 'store',
},
],
});
pass.setPipeline(cellPipeline);
pass.setBindGroup(0, bindGroups[step % 2]);
pass.setVertexBuffer(0, vertexBuffer);
pass.draw(vertices.length / 2, GRID_SIZE * GRID_SIZE);
pass.end();
device.queue.submit([encoder.finish()]);
};
setInterval(updateGrid, UPDATE_INTERVAL);
};
const init = () => {
const canvas = document.querySelector('#canvas');
canvas.width = 300;
canvas.height = 300;
if (!canvas) {
console.log('no canvas');
return;
}
const webgpu = canvas.getContext('webgpu');
if (!webgpu) {
console.log('no webgpu context');
return;
}
if (!navigator.gpu) {
console.log('no navigator.gpu');
return;
}
navigator.gpu.requestAdapter()
.then((adapter) => {
adapter.requestDevice()
.then((device) => {
main(canvas, webgpu, adapter, device);
})
.catch((error) => {
console.log('no gpu device');
console.log(error);
});
})
.catch((error) => {
console.log('no gpu adapter');
console.log(error);
});
};
init();
</script>
</body>
</html>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment