wip

Author: Sheraff

src/pages/particle-life-gpu/bin-fill.wgsl

@@ -0,0 +1,78 @@
+struct Config {
+	binWidth: f32,
+	binHeight: f32,
+	widthDivisions: u32,
+	binCount: u32,
+	particleCount: u32,
+}
+
+@group(0) @binding(0) var<uniform> config: Config;
+@group(1) @binding(0) var<storage, read_write> binSize: array<atomic<u32>>;
+@group(1) @binding(1) var<storage, read_write> binOffset: array<u32>;
+@group(1) @binding(2) var<storage, read_write> binCursor: array<atomic<u32>>;
+@group(1) @binding(3) var<storage, read_write> binContents: array<u32>;
+@group(2) @binding(0) var<storage, read> particlePositions: array<vec2f>;
+
+@compute @workgroup_size(256)
+fn clear(
+	@builtin(global_invocation_id) id: vec3u,
+) {
+	let i = id.x;
+	if (i >= config.binCount) {
+		return;
+	}
+	atomicStore(&binSize[i], 0u);
+	binOffset[i] = 0u;
+	atomicStore(&binCursor[i], 0u);
+}
+
+@compute @workgroup_size(256)
+fn size(
+	@builtin(global_invocation_id) id: vec3u,
+) {
+	let i = id.x;
+	if (i >= config.particleCount) {
+		return;
+	}
+
+	let position = particlePositions[i];
+	let bin_x = u32(position.x / config.binWidth);
+	let bin_y = u32(position.y / config.binHeight);
+	let index = bin_y * config.widthDivisions + bin_x;
+
+	atomicAdd(&binSize[index], 1u);
+}
+
+@compute @workgroup_size(256)
+fn prepare(
+	@builtin(global_invocation_id) id: vec3u,
+) {
+	let i = id.x;
+	if (i >= config.binCount) {
+		return;
+	}
+	var offset = 0u;
+	for (var j: u32 = 0u; j < i; j = j + 1u) {
+		offset = offset + atomicLoad(&binSize[j]);
+	}
+	binOffset[i] = offset;
+}
+
+@compute @workgroup_size(256)
+fn fill(
+	@builtin(global_invocation_id) id: vec3u,
+) {
+	let i = id.x;
+	if (i >= config.particleCount) {
+		return;
+	}
+
+	let position = particlePositions[i];
+	let bin_x = u32(position.x / config.binWidth);
+	let bin_y = u32(position.y / config.binHeight);
+	let index = bin_y * config.widthDivisions + bin_x;
+
+	let offset = atomicAdd(&binCursor[index], 1u);
+	let writeIndex = binOffset[index] + offset;
+	binContents[writeIndex] = i;
+}

src/pages/particle-life-gpu/compute.wgsl

@@ -1,5 +1,5 @@
- @group(0) @binding(0) var<storage, read_write> data: array<f32>;
- 
+@group(0) @binding(0) var<storage, read_write> data: array<f32>;
+
 @compute @workgroup_size(1) fn computeSomething(
 	@builtin(global_invocation_id) id: vec3u
 ) {

src/pages/particle-life-gpu/draw.wgsl

@@ -0,0 +1,56 @@
+struct Config {
+	width: f32,
+	height: f32,
+}
+
+@group(0) @binding(0) var<uniform> config: Config;
+@group(1) @binding(0) var<storage, read> particlePositions: array<vec2f>;
+@group(1) @binding(1) var<storage, read> particleColors: array<u32>;
+
+const size = 1.0;
+const positions = array<vec2f, 6>(
+	vec2f(-1, -1),
+	vec2f( 1, -1),
+	vec2f( 1,  1),
+	vec2f(-1, -1),
+	vec2f( 1,  1),
+	vec2f(-1,  1),
+);
+
+const colors = array<vec4f, 6>(
+	vec4f(1.0, 0.0, 0.0, 1.0), // red
+	vec4f(0.0, 1.0, 0.0, 1.0), // green
+	vec4f(0.0, 0.0, 1.0, 1.0), // blue
+	vec4f(1.0, 1.0, 0.0, 1.0), // yellow
+	vec4f(1.0, 0.0, 1.0, 1.0), // magenta
+	vec4f(0.0, 1.0, 1.0, 1.0), // cyan
+);
+
+struct VertexOutput {
+	// corner of the square (clip space)
+	@builtin(position) position: vec4f,
+	// color of the particle
+	@location(1) color: vec4f,
+}
+
+
+// for each particle, draw a square centered at its position (i.e. 6 vertices per particle to make a square)
+@vertex fn vs(
+	@builtin(instance_index) instanceID: u32
+) -> VertexOutput {
+	let pos = particlePositions[instanceID]; // screen space
+	let vertex_pos = pos; // screen space
+	let clip_vertex_pos = vec4f(
+		(vertex_pos.x / (config.width / 2.0)) - 1.0,
+		-((vertex_pos.y / (config.height / 2.0)) - 1.0),
+		0.0,
+		1.0
+	);
+	let color = colors[particleColors[instanceID]];
+	return VertexOutput(clip_vertex_pos, color);
+}
+
+// paint a circle centered at offset
+@fragment fn fs(data: VertexOutput) -> @location(0) vec4f {
+	return data.color;
+}