renderable = {
const [w, h] = [width, 640]
const { device, context, format } = await gpu.init(w, h)
const PARTICLE_COUNT = 100_000
const PARTICLE_SIZE = 5
const WORKGROUP_SIZE = 64
const WORKGROUP_DISPATCH_COUNT = Math.ceil(PARTICLE_COUNT / WORKGROUP_SIZE)
const module = device.createShaderModule({
label: 'vert & frag shader module',
code: `
struct Uniforms {
size: f32,
elapsed: f32,
mouse: vec2f,
modelViewProjectionMatrix: mat4x4f,
resolution: vec2f,
}
struct Agent {
position: vec3f,
pad1: u32,
heading: vec2f,
pad2: vec2f,
}
struct VertexOut {
@builtin(position) position: vec4f,
@location(1) norm_index: f32,
@location(2) pos: vec3f,
@location(3) heading: vec2f,
};
@group(0) @binding(0) var<uniform> uniforms: Uniforms;
@group(0) @binding(1) var<storage, read> particles: array<Agent>;
@vertex
fn vs(
@builtin(instance_index) instance_index : u32,
@builtin(vertex_index) vertex_index : u32,
) -> VertexOut {
let quad = 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 position = particles[instance_index].position;
let clip_pos = uniforms.modelViewProjectionMatrix * vec4f(position, 1);
let point_pos = vec4f((quad[vertex_index] - 0.5) * (uniforms.size / uniforms.resolution), 0, 0);
let out_pos = clip_pos + point_pos;
return VertexOut(
out_pos,
f32(instance_index) / f32(arrayLength(&particles)),
position,
particles[instance_index].heading,
);
}
@fragment
fn fs(vout: VertexOut) -> @location(0) vec4f {
return vec4f(vout.pos.x * 0.5 + 0.5, vout.pos.y * 0.5 + 0.5, vout.pos.z * 0.5 + 0.5, 1.0) * 0.5;
}
`,
});
const computeShader = `
struct Uniforms {
size: f32,
elapsed: f32,
mouse: vec2f,
modelViewProjectionMatrix: mat4x4f,
resolution: vec2f,
}
struct Agent {
position: vec3f,
pad1: u32,
heading: vec2f,
pad2: vec2f,
}
${ShaderFragment.hash}
const S: vec3f = vec3f(${S});
// pos from -1 to 1 -> 0 to 1 -> 0 to S.* -> index
fn get_index_for_pos(pos: vec3f) -> u32 {
return u32(
(pos.x * 0.5 + 0.5) * S.x +
(pos.y * 0.5 + 0.5) * S.y * S.x +
(pos.z * 0.5 + 0.5) * S.z * (S.x * S.y)
);
}
@group(0) @binding(0) var<uniform> uniforms: Uniforms;
@group(0) @binding(1) var<storage> particles_in: array<Agent>;
@group(0) @binding(2) var<storage, read_write> particles_out: array<Agent>;
@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;
}
let agent = particles_in[index];
particles_out[index].position.y -= 0.01;
if (particles_out[index].position.y <= -1) {
particles_out[index].position.y += 2;
}
}
`;
const computeModule = device.createShaderModule({
label: 'compute shader module',
code: computeShader
})
const bindGroupLayout = device.createBindGroupLayout({
label: 'bind group layout',
entries: [
{
binding: 0,
visibility: GPUShaderStage.VERTEX | GPUShaderStage.COMPUTE,
buffer: {}
},
{
binding: 1,
visibility: GPUShaderStage.VERTEX | GPUShaderStage.COMPUTE,
buffer: { type: 'read-only-storage' }
},
{
binding: 2,
visibility: GPUShaderStage.COMPUTE,
buffer: { type: 'storage' }
},
]
})
const uniforms = new Float32Array([
PARTICLE_SIZE,
0,
0, 0,
...util.arr(16, 0),
w, h,
0, 0
])
const uniformBuffer = device.createBuffer({
label: 'uniforms buffer',
size: uniforms.byteLength,
usage: GPUBufferUsage.UNIFORM | GPUBufferUsage.COPY_DST,
})
device.queue.writeBuffer(uniformBuffer, 0, uniforms)
const matrixValue = uniforms.subarray(4, 4 + 16)
const points = new Float32Array(
util.flatArr(PARTICLE_COUNT, () => [
util.rands(), util.rands(), util.rands(),
0,
util.rands(), util.rands(),
0, 0,
])
)
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: gpu.Blending.Additive,
}],
},
})
const computePipeline = device.createComputePipeline({
label: 'compute pipeline',
layout: pipelineLayout,
compute: {
module: computeModule,
entryPoint: 'cs'
}
})
let step = 0;
let last = 0;
let elapsed;
const projection = mat4.perspective(90 * Math.PI / 180, w / h, 0.1, 50);
const view = mat4.lookAt(
[0, 0, 2],
[0, 0, 0],
[0, 1, 0],
);
const viewProjection = mat4.multiply(projection, view);
mat4.copy(viewProjection, matrixValue)
function render() {
let t = performance.now() * 0.0001
elapsed = t - last
last = t
uniforms.set([elapsed], 1);
mat4.rotateY(viewProjection, t, matrixValue);
device.queue.writeBuffer(uniformBuffer, 0, uniforms);
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()]);
}
return { context, render }
}