From a7727c6b5421195c4b7d3106becd4609bfee9636 Mon Sep 17 00:00:00 2001 From: Andrew Gutekanst Date: Tue, 12 Apr 2022 21:15:49 -0400 Subject: [PATCH] examples: add ported boids example Ported from https://github.com/austinEng/webgpu-samples/ --- build.zig | 27 ++-- examples/boids/main.zig | 220 ++++++++++++++++++++++++++++++ examples/boids/sprite.wgsl | 15 ++ examples/boids/updateSprites.wgsl | 86 ++++++++++++ 4 files changed, 336 insertions(+), 12 deletions(-) create mode 100644 examples/boids/main.zig create mode 100644 examples/boids/sprite.wgsl create mode 100644 examples/boids/updateSprites.wgsl diff --git a/build.zig b/build.zig index 2b47d078..70b28eb8 100644 --- a/build.zig +++ b/build.zig @@ -23,19 +23,22 @@ pub fn build(b: *std.build.Builder) void { const test_step = b.step("test", "Run library tests"); test_step.dependOn(&main_tests.step); - const example = b.addExecutable("hello-triangle", "examples/triangle/main.zig"); - example.setTarget(target); - example.setBuildMode(mode); - example.addPackage(pkg); - example.addPackage(gpu.pkg); - example.addPackage(glfw.pkg); - link(b, example, options); - example.install(); - const example_run_cmd = example.run(); - example_run_cmd.step.dependOn(b.getInstallStep()); - const example_run_step = b.step("run-example", "Run the example"); - example_run_step.dependOn(&example_run_cmd.step); + inline for ([_][] const u8{"triangle", "boids"}) |name| { + const example = b.addExecutable("example-" ++ name, "examples/" ++ name ++ "/main.zig"); + example.setTarget(target); + example.setBuildMode(mode); + example.addPackage(pkg); + example.addPackage(gpu.pkg); + example.addPackage(glfw.pkg); + link(b, example, options); + example.install(); + + const example_run_cmd = example.run(); + example_run_cmd.step.dependOn(b.getInstallStep()); + const example_run_step = b.step("run-example-" ++ name, "Run the example"); + example_run_step.dependOn(&example_run_cmd.step); + } } pub const Options = struct { diff --git a/examples/boids/main.zig b/examples/boids/main.zig new file mode 100644 index 00000000..1e1e31e0 --- /dev/null +++ b/examples/boids/main.zig @@ -0,0 +1,220 @@ +/// A port of Austin Eng's "computeBoids" webgpu sample. +/// https://github.com/austinEng/webgpu-samples/blob/main/src/sample/computeBoids/main.ts + +const std = @import("std"); +const mach = @import("mach"); +const gpu = @import("gpu"); + +const FrameParams = struct { + compute_pipeline: gpu.ComputePipeline, + render_pipeline: gpu.RenderPipeline, + sprite_vertex_buffer: gpu.Buffer, + particle_buffers: [2]gpu.Buffer, + particle_bind_groups: [2]gpu.BindGroup, + frame_counter: usize, +}; +const App = mach.App(*FrameParams, .{}); + +const num_particle = 1500; + +pub fn main() !void { + var gpa = std.heap.GeneralPurposeAllocator(.{}){}; + var allocator = gpa.allocator(); + + const ctx = try allocator.create(FrameParams); + var app = try App.init(allocator, ctx, .{}); + + const sprite_shader_module = app.device.createShaderModule(&.{ + .label = "sprite shader module", + .code = .{ .wgsl = @embedFile("sprite.wgsl") }, + }); + + const update_sprite_shader_module = app.device.createShaderModule(&.{ + .label = "update sprite shader module", + .code = .{ .wgsl = @embedFile("updateSprites.wgsl") }, + }); + + const instanced_particles_attributes = [_]gpu.VertexAttribute{ + .{ + // instance position + .shader_location = 0, + .offset = 0, + .format = .float32x2, + }, + .{ + // instance velocity + .shader_location = 1, + .offset = 2 * 4, + .format = .float32x2, + }, + }; + + const vertex_buffer_attributes = [_]gpu.VertexAttribute{ + .{ + // vertex positions + .shader_location = 2, + .offset = 0, + .format = .float32x2, + }, + }; + + const render_pipeline = app.device.createRenderPipeline(&gpu.RenderPipeline.Descriptor{ + .vertex = .{ + .module = sprite_shader_module, + .entry_point = "vert_main", + .buffers = &[_]gpu.VertexBufferLayout{ + .{ + // instanced particles buffer + .array_stride = 4*4, + .step_mode = .instance, + .attribute_count = instanced_particles_attributes.len, + .attributes = &instanced_particles_attributes, + }, + .{ + // vertex buffer + .array_stride = 2*4, + .step_mode = .vertex, + .attribute_count = vertex_buffer_attributes.len, + .attributes = &vertex_buffer_attributes, + }, + }, + }, + .fragment = &gpu.FragmentState{ + .module = sprite_shader_module, + .entry_point = "frag_main", + .targets = &[_]gpu.ColorTargetState{ + .{ + .format = app.swap_chain_format, + }, + } + }, + }); + + const compute_pipeline = app.device.createComputePipeline(&gpu.ComputePipeline.Descriptor{ + .compute = gpu.ProgrammableStageDescriptor{ + .module = update_sprite_shader_module, + .entry_point = "main", + } + }); + + const vert_buffer_data = [_]f32{ + -0.01, -0.02, 0.01, + -0.02, 0.0, 0.02, + }; + + const sprite_vertex_buffer = app.device.createBuffer(&gpu.Buffer.Descriptor{ + .usage = .{.vertex = true, .copy_dst = true}, + .size = vert_buffer_data.len * @sizeOf(f32), + }); + app.device.getQueue().writeBuffer(sprite_vertex_buffer, 0, f32, &vert_buffer_data); + + const sim_params = [_]f32 { + 0.04, // .delta_T + 0.1, // .rule_1_distance + 0.025, // .rule_2_distance + 0.025, // .rule_3_distance + 0.02, // .rule_1_scale + 0.05, // .rule_2_scale + 0.005, // .rule_3_scale + }; + + const sim_param_buffer = app.device.createBuffer(&gpu.Buffer.Descriptor{ + .usage = .{.uniform = true, .copy_dst = true}, + .size = sim_params.len * @sizeOf(f32), + }); + app.device.getQueue().writeBuffer(sim_param_buffer, 0, f32, &sim_params); + + var initial_particle_data: [num_particle*4]f32 = undefined; + var rng = std.rand.DefaultPrng.init(0); + const random = rng.random(); + var i:usize = 0; + while(i < num_particle): (i += 1) { + initial_particle_data[4 * i + 0] = 2 * (random.float(f32) - 0.5); + initial_particle_data[4 * i + 1] = 2 * (random.float(f32) - 0.5); + initial_particle_data[4 * i + 2] = 2 * (random.float(f32) - 0.5) * 0.1; + initial_particle_data[4 * i + 3] = 2 * (random.float(f32) - 0.5) * 0.1; + } + + var particle_buffers: [2]gpu.Buffer = undefined; + var particle_bind_groups: [2]gpu.BindGroup = undefined; + i = 0; + while(i < 2): (i+=1) { + particle_buffers[i] = app.device.createBuffer(&gpu.Buffer.Descriptor{ + .usage = .{.vertex = true, .copy_dst = true, .storage = true, }, + .size = initial_particle_data.len * @sizeOf(f32), + }); + app.device.getQueue().writeBuffer(particle_buffers[i], 0, f32, &initial_particle_data); + } + + i = 0; + while(i < 2): (i+=1) { + particle_bind_groups[i] = app.device.createBindGroup(&gpu.BindGroup.Descriptor{ + .layout = compute_pipeline.getBindGroupLayout(0), + .entries = &[_]gpu.BindGroup.Entry { + gpu.BindGroup.Entry.buffer(0, sim_param_buffer, 0, sim_params.len * @sizeOf(f32)), + gpu.BindGroup.Entry.buffer(1, particle_buffers[i], 0, initial_particle_data.len * @sizeOf(f32)), + gpu.BindGroup.Entry.buffer(2, particle_buffers[(i+1)%2], 0, initial_particle_data.len * @sizeOf(f32)), + } + }); + } + + ctx.* = FrameParams{ + .compute_pipeline = compute_pipeline, + .render_pipeline = render_pipeline, + .sprite_vertex_buffer = sprite_vertex_buffer, + .particle_buffers = particle_buffers, + .particle_bind_groups = particle_bind_groups, + .frame_counter = 0, + }; + + try app.run(.{ .frame = frame }); +} + +fn frame(app: *App, params: *FrameParams) !void { + const back_buffer_view = app.swap_chain.?.getCurrentTextureView(); + const color_attachment = gpu.RenderPassColorAttachment{ + .view = back_buffer_view, + .resolve_target = null, + .clear_value = std.mem.zeroes(gpu.Color), + .load_op = .clear, + .store_op = .store, + }; + + const render_pass_descriptor = gpu.RenderPassEncoder.Descriptor{ + .color_attachments = &[_]gpu.RenderPassColorAttachment { + color_attachment, + } + }; + + const command_encoder = app.device.createCommandEncoder(null); + { + const pass_encoder = command_encoder.beginComputePass(null); + pass_encoder.setPipeline(params.compute_pipeline); + pass_encoder.setBindGroup(0, params.particle_bind_groups[params.frame_counter % 2], null); + pass_encoder.dispatch(@floatToInt(u32, std.math.ceil(@as(f32, num_particle) / 64)), 1, 1); + pass_encoder.end(); + pass_encoder.release(); + } + { + const pass_encoder = command_encoder.beginRenderPass(&render_pass_descriptor); + pass_encoder.setPipeline(params.render_pipeline); + pass_encoder.setVertexBuffer(0, params.particle_buffers[(params.frame_counter + 1) % 2], 0, num_particle*4*@sizeOf(f32)); + pass_encoder.setVertexBuffer(1, params.sprite_vertex_buffer, 0, 6*@sizeOf(f32)); + pass_encoder.draw(3, num_particle, 0, 0); + pass_encoder.end(); + pass_encoder.release(); + } + + params.frame_counter += 1; + if(params.frame_counter % 60 == 0) { + std.debug.print("Frame {}\n", .{params.frame_counter}); + } + + var command = command_encoder.finish(null); + command_encoder.release(); + app.device.getQueue().submit(&.{command}); + command.release(); + + app.swap_chain.?.present(); + back_buffer_view.release(); +} \ No newline at end of file diff --git a/examples/boids/sprite.wgsl b/examples/boids/sprite.wgsl new file mode 100644 index 00000000..63a6c506 --- /dev/null +++ b/examples/boids/sprite.wgsl @@ -0,0 +1,15 @@ +@stage(vertex) +fn vert_main(@location(0) a_particlePos : vec2, + @location(1) a_particleVel : vec2, + @location(2) a_pos : vec2) -> @builtin(position) vec4 { + let angle = -atan2(a_particleVel.x, a_particleVel.y); + let pos = vec2( + (a_pos.x * cos(angle)) - (a_pos.y * sin(angle)), + (a_pos.x * sin(angle)) + (a_pos.y * cos(angle))); + return vec4(pos + a_particlePos, 0.0, 1.0); +} + +@stage(fragment) +fn frag_main() -> @location(0) vec4 { + return vec4(1.0, 1.0, 1.0, 1.0); +} diff --git a/examples/boids/updateSprites.wgsl b/examples/boids/updateSprites.wgsl new file mode 100644 index 00000000..cefba51a --- /dev/null +++ b/examples/boids/updateSprites.wgsl @@ -0,0 +1,86 @@ +struct Particle { + pos : vec2; + vel : vec2; +}; +struct SimParams { + deltaT : f32; + rule1Distance : f32; + rule2Distance : f32; + rule3Distance : f32; + rule1Scale : f32; + rule2Scale : f32; + rule3Scale : f32; +}; +struct Particles { + particles : array; +}; +@binding(0) @group(0) var params : SimParams; +@binding(1) @group(0) var particlesA : Particles; +@binding(2) @group(0) var particlesB : Particles; + +// https://github.com/austinEng/Project6-Vulkan-Flocking/blob/master/data/shaders/computeparticles/particle.comp +@stage(compute) @workgroup_size(64) +fn main(@builtin(global_invocation_id) GlobalInvocationID : vec3) { + var index : u32 = GlobalInvocationID.x; + + var vPos = particlesA.particles[index].pos; + var vVel = particlesA.particles[index].vel; + var cMass = vec2(0.0, 0.0); + var cVel = vec2(0.0, 0.0); + var colVel = vec2(0.0, 0.0); + var cMassCount : u32 = 0u; + var cVelCount : u32 = 0u; + var pos : vec2; + var vel : vec2; + + for (var i : u32 = 0u; i < arrayLength(&particlesA.particles); i = i + 1u) { + if (i == index) { + continue; + } + + pos = particlesA.particles[i].pos.xy; + vel = particlesA.particles[i].vel.xy; + if (distance(pos, vPos) < params.rule1Distance) { + cMass = cMass + pos; + cMassCount = cMassCount + 1u; + } + if (distance(pos, vPos) < params.rule2Distance) { + colVel = colVel - (pos - vPos); + } + if (distance(pos, vPos) < params.rule3Distance) { + cVel = cVel + vel; + cVelCount = cVelCount + 1u; + } + } + if (cMassCount > 0u) { + var temp = f32(cMassCount); + cMass = (cMass / vec2(temp, temp)) - vPos; + } + if (cVelCount > 0u) { + var temp = f32(cVelCount); + cVel = cVel / vec2(temp, temp); + } + vVel = vVel + (cMass * params.rule1Scale) + (colVel * params.rule2Scale) + + (cVel * params.rule3Scale); + + // clamp velocity for a more pleasing simulation + vVel = normalize(vVel) * clamp(length(vVel), 0.0, 0.1); + // kinematic update + vPos = vPos + (vVel * params.deltaT); + // Wrap around boundary + if (vPos.x < -1.0) { + vPos.x = 1.0; + } + if (vPos.x > 1.0) { + vPos.x = -1.0; + } + if (vPos.y < -1.0) { + vPos.y = 1.0; + } + if (vPos.y > 1.0) { + vPos.y = -1.0; + } + // Write back + particlesB.particles[index].pos = vPos; + particlesB.particles[index].vel = vVel; +}