renderable = {
const { device, context, format } = await gpu.init(768, 768)
const PARTICLE_COUNT = 1_000
const PARTICLE_SIZE = 0.01
const WORKGROUP_SIZE = 64
const WORKGROUP_DISPATCH_COUNT = Math.ceil(PARTICLE_COUNT / WORKGROUP_SIZE)
const module = device.createShaderModule({
label: 'vert & frag shader module',
code: `
${shaderFragments.hsl2rgb}
struct Uniforms {
size: f32,
count: f32,
mouse: vec2f,
}
struct Particle {
position: vec2f,
velocity: vec2f,
}
struct VertexOut {
@builtin(position) position: vec4f,
@location(1) norm_index: f32,
@location(2) pos: vec2f,
};
@group(0) @binding(0) var<uniform> uniforms: Uniforms;
@group(0) @binding(1) var<storage, read> particles: array<Particle>;
@vertex
fn vs(
@builtin(instance_index) instance_index : u32,
@builtin(vertex_index) vertex_index : u32,
) -> VertexOut {
let p = array(
vec2f(0.0, 0.0),
vec2f(1.0, 0.0),
vec2f(0.0, 1.0),
vec2f(0.0, 1.0),
vec2f(1.0, 0.0),
vec2f(1.0, 1.0),
);
let pos = particles[instance_index].position;
let xy = (p[vertex_index] - 0.5) * uniforms.size + pos;
return VertexOut(
vec4f(xy, 1, 1),
f32(instance_index) / f32(arrayLength(&particles)),
pos
);
}
@fragment
fn fs(vout: VertexOut) -> @location(0) vec4f {
return vec4f(hsl2rgb(vec3f(vout.norm_index, 1, 0.65)), 1);
}
`,
});
const computeShader = `
struct Uniforms {
size: f32,
count: f32,
mouse: vec2f,
}
struct Particle {
position: vec2f,
velocity: vec2f,
}
@group(0) @binding(0) var<uniform> uniforms: Uniforms;
@group(0) @binding(1) var<storage> particles_in: array<Particle>;
@group(0) @binding(2) var<storage, read_write> particles_out: array<Particle>;
@compute
@workgroup_size(${WORKGROUP_SIZE})
fn cs(@builtin(global_invocation_id) global_invocation_id: vec3u) {
let index = global_invocation_id.x;
if (index > arrayLength(&particles_in)) {
return;
}
var next_x = particles_in[index].position.x + 0.01;
// there's a more efficient way to do this...
if (next_x > 1) {
next_x = -1;
}
particles_out[index].position.x = next_x;
}
`
const computeModule = device.createShaderModule({
label: 'compute shader module',
code: computeShader
})
// define access to resources across all pipelines
const bindGroupLayout = device.createBindGroupLayout({
label: 'bind group layout',
entries: [
// uniforms
{
binding: 0,
visibility: GPUShaderStage.VERTEX | GPUShaderStage.COMPUTE,
buffer: {}
},
// particles in
{
binding: 1,
visibility: GPUShaderStage.VERTEX | GPUShaderStage.COMPUTE,
buffer: { type: 'read-only-storage' }
},
// particles out
{
binding: 2,
visibility: GPUShaderStage.COMPUTE,
buffer: { type: 'storage' }
},
]
})
// note about alignment...
// (https://surma.dev/things/webgpu/ ctrl-f "alignment")
const uniforms = new Float32Array([
// this works because the vec2f mouse coords must be aligned to a memory
// address of 8; particle size (f32, size=4, align=4) + particle count (f32) = 8
// which allows us to align the size=8 vec2<f32> align=8 to a memory address
// multiple of 8
PARTICLE_SIZE, // 4
PARTICLE_COUNT, // + 4
0, 0 // = address 8
])
const uniformBuffer = device.createBuffer({
label: 'uniforms buffer',
size: uniforms.byteLength,
usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST,
})
device.queue.writeBuffer(uniformBuffer, 0, uniforms);
// create subarray to avoid byte offsets later
const uMouse = uniforms.subarray(4 + 4, 4 + 4 + 4 * 2)
// position x, position y, velocity x, velocity y
const points = new Float32Array(util.arr(PARTICLE_COUNT * 4, () => util.randn(0, 0.1)).flat())
const pointsBufferA = device.createBuffer({
label: 'points storage buffer A',
size: points.byteLength,
usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST
})
device.queue.writeBuffer(pointsBufferA, 0, points);
const pointsBufferB = device.createBuffer({
label: 'points storage buffer B',
size: points.byteLength,
usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST
})
device.queue.writeBuffer(pointsBufferB, 0, points);
const bindGroupA = device.createBindGroup({
label: 'bind group layout A',
layout: bindGroupLayout,
entries: [
{ binding: 0, resource: { buffer: uniformBuffer }},
{ binding: 1, resource: { buffer: pointsBufferA }},
{ binding: 2, resource: { buffer: pointsBufferB }},
]
});
const bindGroupB = device.createBindGroup({
label: 'bind group layout B',
layout: bindGroupLayout,
entries: [
{ binding: 0, resource: { buffer: uniformBuffer }},
{ binding: 1, resource: { buffer: pointsBufferB }},
{ binding: 2, resource: { buffer: pointsBufferA }},
]
});
const pingPong = [bindGroupA, bindGroupB]
const pipelineLayout = device.createPipelineLayout({
label: 'pipeline layout',
bindGroupLayouts: [bindGroupLayout]
})
const pipeline = device.createRenderPipeline({
label: 'pipeline',
layout: pipelineLayout,
vertex: {
module,
entryPoint: 'vs',
},
fragment: {
module,
entryPoint: 'fs',
targets: [{
format,
blend: {
color: {
srcFactor: 'one',
dstFactor: 'one-minus-src-alpha'
},
alpha: {
srcFactor: 'one',
dstFactor: 'one-minus-src-alpha'
},
},
}],
},
})
const computePipeline = device.createComputePipeline({
label: 'compute pipeline',
layout: pipelineLayout,
compute: {
module: computeModule,
entryPoint: 'cs'
}
})
let step = 0;
function render() {
const encoder = device.createCommandEncoder()
const computePass = encoder.beginComputePass()
computePass.setPipeline(computePipeline)
computePass.setBindGroup(0, pingPong[step % 2])
computePass.dispatchWorkgroups(WORKGROUP_DISPATCH_COUNT)
computePass.end()
step++;
const pass = encoder.beginRenderPass({
colorAttachments: [
{
clearValue: [0, 0, 0, 1],
loadOp: 'clear',
storeOp: 'store',
view: context.getCurrentTexture().createView()
},
],
})
pass.setPipeline(pipeline);
pass.setBindGroup(0, pingPong[step % 2]);
pass.draw(6, PARTICLE_COUNT);
pass.end();
device.queue.submit([encoder.finish()]);
}
render()
return { context, render }
}