Created
January 4, 2024 17:55
-
-
Save shogonir/c5ee2383705f942c4f64504597b18a8f to your computer and use it in GitHub Desktop.
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
<!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