blob: 44183487575f052711ffdce8fdb53bd49b7ea1de [file] [log] [blame]
Corentin Wallez4a9ef4e2018-07-18 11:40:26 +02001// Copyright 2017 The Dawn Authors
Corentin Wallezf07e3bd2017-04-20 14:38:20 -04002//
3// Licensed under the Apache License, Version 2.0 (the "License");
4// you may not use this file except in compliance with the License.
5// You may obtain a copy of the License at
6//
7// http://www.apache.org/licenses/LICENSE-2.0
8//
9// Unless required by applicable law or agreed to in writing, software
10// distributed under the License is distributed on an "AS IS" BASIS,
11// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
12// See the License for the specific language governing permissions and
13// limitations under the License.
14
Corentin Wallez9347e8f2017-06-19 13:15:13 -040015#include "SampleUtils.h"
Corentin Wallezf07e3bd2017-04-20 14:38:20 -040016
Yan, Shaoboa4924272018-12-10 19:47:22 +000017#include "utils/ComboRenderPipelineDescriptor.h"
Austin Eng700a5fb2021-06-24 19:21:31 +000018#include "utils/ScopedAutoreleasePool.h"
Corentin Wallez134e0802017-07-17 17:13:57 -040019#include "utils/SystemUtils.h"
Corentin Wallez04863c42019-10-25 11:36:47 +000020#include "utils/WGPUHelpers.h"
Corentin Wallez5ee7afd2017-06-19 13:09:41 -040021
Corentin Wallezf07e3bd2017-04-20 14:38:20 -040022#include <array>
23#include <cstring>
24#include <random>
Corentin Wallezf07e3bd2017-04-20 14:38:20 -040025
Corentin Wallez04863c42019-10-25 11:36:47 +000026wgpu::Device device;
27wgpu::Queue queue;
28wgpu::SwapChain swapchain;
29wgpu::TextureView depthStencilView;
Corentin Wallezf07e3bd2017-04-20 14:38:20 -040030
Corentin Wallez04863c42019-10-25 11:36:47 +000031wgpu::Buffer modelBuffer;
32std::array<wgpu::Buffer, 2> particleBuffers;
Corentin Wallezf07e3bd2017-04-20 14:38:20 -040033
Corentin Wallez04863c42019-10-25 11:36:47 +000034wgpu::RenderPipeline renderPipeline;
Corentin Wallezf07e3bd2017-04-20 14:38:20 -040035
Corentin Wallez04863c42019-10-25 11:36:47 +000036wgpu::Buffer updateParams;
37wgpu::ComputePipeline updatePipeline;
38std::array<wgpu::BindGroup, 2> updateBGs;
Corentin Wallezf07e3bd2017-04-20 14:38:20 -040039
Corentin Wallezf07e3bd2017-04-20 14:38:20 -040040size_t pingpong = 0;
41
42static const uint32_t kNumParticles = 1000;
43
44struct Particle {
Corentin Wallez5bd94452021-12-14 08:42:36 +000045 std::array<float, 2> pos;
46 std::array<float, 2> vel;
Corentin Wallezf07e3bd2017-04-20 14:38:20 -040047};
48
49struct SimParams {
50 float deltaT;
51 float rule1Distance;
52 float rule2Distance;
53 float rule3Distance;
54 float rule1Scale;
55 float rule2Scale;
56 float rule3Scale;
57 int particleCount;
58};
59
60void initBuffers() {
Corentin Wallez5bd94452021-12-14 08:42:36 +000061 std::array<std::array<float, 2>, 3> model = {{
Corentin Wallezf07e3bd2017-04-20 14:38:20 -040062 {-0.01, -0.02},
63 {0.01, -0.02},
64 {0.00, 0.02},
Corentin Wallez5bd94452021-12-14 08:42:36 +000065 }};
Corentin Wallez9e9e29f2019-08-27 08:21:39 +000066 modelBuffer =
Corentin Wallez5bd94452021-12-14 08:42:36 +000067 utils::CreateBufferFromData(device, &model, sizeof(model), wgpu::BufferUsage::Vertex);
Corentin Wallezf07e3bd2017-04-20 14:38:20 -040068
Kai Ninomiya2afea0c2020-07-10 20:33:08 +000069 SimParams params = {0.04f, 0.1f, 0.025f, 0.025f, 0.02f, 0.05f, 0.005f, kNumParticles};
Corentin Wallez9e9e29f2019-08-27 08:21:39 +000070 updateParams =
Corentin Wallez04863c42019-10-25 11:36:47 +000071 utils::CreateBufferFromData(device, &params, sizeof(params), wgpu::BufferUsage::Uniform);
Corentin Wallezf07e3bd2017-04-20 14:38:20 -040072
73 std::vector<Particle> initialParticles(kNumParticles);
74 {
75 std::mt19937 generator;
76 std::uniform_real_distribution<float> dist(-1.0f, 1.0f);
Kai Ninomiya2afea0c2020-07-10 20:33:08 +000077 for (auto& p : initialParticles) {
Corentin Wallez5bd94452021-12-14 08:42:36 +000078 p.pos = {dist(generator), dist(generator)};
79 p.vel = {dist(generator) * 0.1f, dist(generator) * 0.1f};
Corentin Wallezf07e3bd2017-04-20 14:38:20 -040080 }
81 }
82
Kai Ninomiya78c8b832017-07-21 17:00:22 -070083 for (size_t i = 0; i < 2; i++) {
Corentin Wallez04863c42019-10-25 11:36:47 +000084 wgpu::BufferDescriptor descriptor;
Corentin Wallez82b65732018-08-22 15:37:29 +020085 descriptor.size = sizeof(Particle) * kNumParticles;
Corentin Wallez9e9e29f2019-08-27 08:21:39 +000086 descriptor.usage =
Corentin Wallez04863c42019-10-25 11:36:47 +000087 wgpu::BufferUsage::CopyDst | wgpu::BufferUsage::Vertex | wgpu::BufferUsage::Storage;
Corentin Wallez82b65732018-08-22 15:37:29 +020088 particleBuffers[i] = device.CreateBuffer(&descriptor);
Corentin Wallezf07e3bd2017-04-20 14:38:20 -040089
Corentin Wallez47a33412020-06-02 09:24:39 +000090 queue.WriteBuffer(particleBuffers[i], 0,
91 reinterpret_cast<uint8_t*>(initialParticles.data()),
92 sizeof(Particle) * kNumParticles);
Corentin Wallezf07e3bd2017-04-20 14:38:20 -040093 }
94}
95
96void initRender() {
Corentin Wallez7aec4ae2021-03-24 15:55:32 +000097 wgpu::ShaderModule vsModule = utils::CreateShaderModule(device, R"(
James Price9e0debd2021-04-13 14:52:44 +000098 struct VertexIn {
99 [[location(0)]] a_particlePos : vec2<f32>;
100 [[location(1)]] a_particleVel : vec2<f32>;
101 [[location(2)]] a_pos : vec2<f32>;
102 };
Corentin Wallez4814bdb2020-11-26 16:39:46 +0000103
104 [[stage(vertex)]]
James Price9e0debd2021-04-13 14:52:44 +0000105 fn main(input : VertexIn) -> [[builtin(position)]] vec4<f32> {
106 var angle : f32 = -atan2(input.a_particleVel.x, input.a_particleVel.y);
Corentin Wallez4814bdb2020-11-26 16:39:46 +0000107 var pos : vec2<f32> = vec2<f32>(
James Price9e0debd2021-04-13 14:52:44 +0000108 (input.a_pos.x * cos(angle)) - (input.a_pos.y * sin(angle)),
109 (input.a_pos.x * sin(angle)) + (input.a_pos.y * cos(angle)));
110 return vec4<f32>(pos + input.a_particlePos, 0.0, 1.0);
Corentin Wallezf07e3bd2017-04-20 14:38:20 -0400111 }
112 )");
113
Corentin Wallez7aec4ae2021-03-24 15:55:32 +0000114 wgpu::ShaderModule fsModule = utils::CreateShaderModule(device, R"(
Corentin Wallez4814bdb2020-11-26 16:39:46 +0000115 [[stage(fragment)]]
James Price9e0debd2021-04-13 14:52:44 +0000116 fn main() -> [[location(0)]] vec4<f32> {
117 return vec4<f32>(1.0, 1.0, 1.0, 1.0);
Corentin Wallezf07e3bd2017-04-20 14:38:20 -0400118 }
119 )");
120
Kai Ninomiyac16a67a2017-07-27 18:30:57 -0700121 depthStencilView = CreateDefaultDepthStencilView(device);
122
Brandon Jones41c87d92021-05-21 05:01:38 +0000123 utils::ComboRenderPipelineDescriptor descriptor;
Yunchao He889d7432019-03-27 18:08:50 +0000124
Corentin Wallez9895c272021-03-18 16:46:58 +0000125 descriptor.vertex.module = vsModule;
126 descriptor.vertex.bufferCount = 2;
127 descriptor.cBuffers[0].arrayStride = sizeof(Particle);
Corentin Wallezff6c9ca62021-07-25 18:40:19 +0000128 descriptor.cBuffers[0].stepMode = wgpu::VertexStepMode::Instance;
Corentin Wallez9895c272021-03-18 16:46:58 +0000129 descriptor.cBuffers[0].attributeCount = 2;
130 descriptor.cAttributes[0].offset = offsetof(Particle, pos);
131 descriptor.cAttributes[0].format = wgpu::VertexFormat::Float32x2;
132 descriptor.cAttributes[1].shaderLocation = 1;
133 descriptor.cAttributes[1].offset = offsetof(Particle, vel);
134 descriptor.cAttributes[1].format = wgpu::VertexFormat::Float32x2;
Corentin Wallez5bd94452021-12-14 08:42:36 +0000135 descriptor.cBuffers[1].arrayStride = 2 * sizeof(float);
Corentin Wallez9895c272021-03-18 16:46:58 +0000136 descriptor.cBuffers[1].attributeCount = 1;
137 descriptor.cBuffers[1].attributes = &descriptor.cAttributes[2];
138 descriptor.cAttributes[2].shaderLocation = 2;
139 descriptor.cAttributes[2].format = wgpu::VertexFormat::Float32x2;
Yan, Shaoboa4924272018-12-10 19:47:22 +0000140
Corentin Wallez9895c272021-03-18 16:46:58 +0000141 descriptor.cFragment.module = fsModule;
142 descriptor.EnableDepthStencil(wgpu::TextureFormat::Depth24PlusStencil8);
143 descriptor.cTargets[0].format = GetPreferredSwapChainTextureFormat();
144
Brandon Jones41c87d92021-05-21 05:01:38 +0000145 renderPipeline = device.CreateRenderPipeline(&descriptor);
Corentin Wallezf07e3bd2017-04-20 14:38:20 -0400146}
147
148void initSim() {
Corentin Wallez7aec4ae2021-03-24 15:55:32 +0000149 wgpu::ShaderModule module = utils::CreateShaderModule(device, R"(
Stephen Whited6754472021-02-23 20:30:39 +0000150 struct Particle {
Ben Claytonc5686842021-03-17 09:48:19 +0000151 pos : vec2<f32>;
152 vel : vec2<f32>;
Corentin Wallezf07e3bd2017-04-20 14:38:20 -0400153 };
James Priced4f8c392021-12-15 13:13:26 +0000154 struct SimParams {
Ben Claytonc5686842021-03-17 09:48:19 +0000155 deltaT : f32;
156 rule1Distance : f32;
157 rule2Distance : f32;
158 rule3Distance : f32;
159 rule1Scale : f32;
160 rule2Scale : f32;
161 rule3Scale : f32;
162 particleCount : u32;
Corentin Wallez4814bdb2020-11-26 16:39:46 +0000163 };
James Priced4f8c392021-12-15 13:13:26 +0000164 struct Particles {
Ben Claytonc5686842021-03-17 09:48:19 +0000165 particles : array<Particle>;
Corentin Wallez4814bdb2020-11-26 16:39:46 +0000166 };
dan sinclair0f9c2d72021-01-19 14:18:51 +0000167 [[binding(0), group(0)]] var<uniform> params : SimParams;
Ben Clayton15eba9a2021-06-08 15:36:44 +0000168 [[binding(1), group(0)]] var<storage, read> particlesA : Particles;
169 [[binding(2), group(0)]] var<storage, read_write> particlesB : Particles;
Corentin Wallezf07e3bd2017-04-20 14:38:20 -0400170
dan sinclaire6ca2542021-01-12 22:11:14 +0000171 // https://github.com/austinEng/Project6-Vulkan-Flocking/blob/master/data/shaders/computeparticles/particle.comp
Sarah98ab91b2021-06-29 00:10:01 +0000172 [[stage(compute), workgroup_size(1)]]
James Price9e0debd2021-04-13 14:52:44 +0000173 fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
Corentin Wallez4814bdb2020-11-26 16:39:46 +0000174 var index : u32 = GlobalInvocationID.x;
175 if (index >= params.particleCount) {
176 return;
177 }
178 var vPos : vec2<f32> = particlesA.particles[index].pos;
179 var vVel : vec2<f32> = particlesA.particles[index].vel;
180 var cMass : vec2<f32> = vec2<f32>(0.0, 0.0);
181 var cVel : vec2<f32> = vec2<f32>(0.0, 0.0);
182 var colVel : vec2<f32> = vec2<f32>(0.0, 0.0);
183 var cMassCount : u32 = 0u;
184 var cVelCount : u32 = 0u;
185 var pos : vec2<f32>;
186 var vel : vec2<f32>;
Corentin Wallezf07e3bd2017-04-20 14:38:20 -0400187
Corentin Wallez4814bdb2020-11-26 16:39:46 +0000188 for (var i : u32 = 0u; i < params.particleCount; i = i + 1u) {
189 if (i == index) {
190 continue;
191 }
Corentin Wallezf07e3bd2017-04-20 14:38:20 -0400192
Kai Ninomiya7883e7e2018-07-10 11:01:28 -0700193 pos = particlesA.particles[i].pos.xy;
194 vel = particlesA.particles[i].vel.xy;
Corentin Wallezf07e3bd2017-04-20 14:38:20 -0400195 if (distance(pos, vPos) < params.rule1Distance) {
Corentin Wallez4814bdb2020-11-26 16:39:46 +0000196 cMass = cMass + pos;
197 cMassCount = cMassCount + 1u;
Corentin Wallezf07e3bd2017-04-20 14:38:20 -0400198 }
199 if (distance(pos, vPos) < params.rule2Distance) {
Corentin Wallez4814bdb2020-11-26 16:39:46 +0000200 colVel = colVel - (pos - vPos);
Corentin Wallezf07e3bd2017-04-20 14:38:20 -0400201 }
202 if (distance(pos, vPos) < params.rule3Distance) {
Corentin Wallez4814bdb2020-11-26 16:39:46 +0000203 cVel = cVel + vel;
204 cVelCount = cVelCount + 1u;
Corentin Wallezf07e3bd2017-04-20 14:38:20 -0400205 }
206 }
Corentin Wallez4814bdb2020-11-26 16:39:46 +0000207
208 if (cMassCount > 0u) {
209 cMass = (cMass / vec2<f32>(f32(cMassCount), f32(cMassCount))) - vPos;
Corentin Wallezf07e3bd2017-04-20 14:38:20 -0400210 }
211
Corentin Wallez4814bdb2020-11-26 16:39:46 +0000212 if (cVelCount > 0u) {
213 cVel = cVel / vec2<f32>(f32(cVelCount), f32(cVelCount));
214 }
215 vVel = vVel + (cMass * params.rule1Scale) + (colVel * params.rule2Scale) +
216 (cVel * params.rule3Scale);
Corentin Wallezf07e3bd2017-04-20 14:38:20 -0400217
dan sinclaire6ca2542021-01-12 22:11:14 +0000218 // clamp velocity for a more pleasing simulation
Corentin Wallezf07e3bd2017-04-20 14:38:20 -0400219 vVel = normalize(vVel) * clamp(length(vVel), 0.0, 0.1);
dan sinclaire6ca2542021-01-12 22:11:14 +0000220 // kinematic update
Corentin Wallez4814bdb2020-11-26 16:39:46 +0000221 vPos = vPos + (vVel * params.deltaT);
Corentin Wallezf07e3bd2017-04-20 14:38:20 -0400222
dan sinclaire6ca2542021-01-12 22:11:14 +0000223 // Wrap around boundary
Corentin Wallez4814bdb2020-11-26 16:39:46 +0000224 if (vPos.x < -1.0) {
225 vPos.x = 1.0;
226 }
227 if (vPos.x > 1.0) {
228 vPos.x = -1.0;
229 }
230 if (vPos.y < -1.0) {
231 vPos.y = 1.0;
232 }
233 if (vPos.y > 1.0) {
234 vPos.y = -1.0;
235 }
Corentin Wallezf07e3bd2017-04-20 14:38:20 -0400236
dan sinclaire6ca2542021-01-12 22:11:14 +0000237 // Write back
Kai Ninomiya7883e7e2018-07-10 11:01:28 -0700238 particlesB.particles[index].pos = vPos;
Kai Ninomiya7883e7e2018-07-10 11:01:28 -0700239 particlesB.particles[index].vel = vVel;
Corentin Wallez4814bdb2020-11-26 16:39:46 +0000240 return;
Corentin Wallezf07e3bd2017-04-20 14:38:20 -0400241 }
242 )");
243
Kai Ninomiya234becf2018-07-10 12:23:50 -0700244 auto bgl = utils::MakeBindGroupLayout(
245 device, {
Corentin Wallez9895c272021-03-18 16:46:58 +0000246 {0, wgpu::ShaderStage::Compute, wgpu::BufferBindingType::Uniform},
247 {1, wgpu::ShaderStage::Compute, wgpu::BufferBindingType::Storage},
248 {2, wgpu::ShaderStage::Compute, wgpu::BufferBindingType::Storage},
Kai Ninomiya234becf2018-07-10 12:23:50 -0700249 });
Corentin Wallezf07e3bd2017-04-20 14:38:20 -0400250
Corentin Wallez04863c42019-10-25 11:36:47 +0000251 wgpu::PipelineLayout pl = utils::MakeBasicPipelineLayout(device, &bgl);
Corentin Wallezf07e3bd2017-04-20 14:38:20 -0400252
Corentin Wallez04863c42019-10-25 11:36:47 +0000253 wgpu::ComputePipelineDescriptor csDesc;
Corentin Wallezaa7109c2018-10-25 10:42:49 +0000254 csDesc.layout = pl;
Brandon Jones0d50a2c2021-06-09 18:07:32 +0000255 csDesc.compute.module = module;
256 csDesc.compute.entryPoint = "main";
Corentin Wallez8e335a52018-08-27 23:12:56 +0200257 updatePipeline = device.CreateComputePipeline(&csDesc);
Corentin Wallezf07e3bd2017-04-20 14:38:20 -0400258
Corentin Wallezf07e3bd2017-04-20 14:38:20 -0400259 for (uint32_t i = 0; i < 2; ++i) {
Kai Ninomiya2afea0c2020-07-10 20:33:08 +0000260 updateBGs[i] = utils::MakeBindGroup(
261 device, bgl,
262 {
263 {0, updateParams, 0, sizeof(SimParams)},
264 {1, particleBuffers[i], 0, kNumParticles * sizeof(Particle)},
265 {2, particleBuffers[(i + 1) % 2], 0, kNumParticles * sizeof(Particle)},
266 });
Corentin Wallezf07e3bd2017-04-20 14:38:20 -0400267 }
268}
269
Corentin Wallez604072b2019-11-12 18:30:11 +0000270wgpu::CommandBuffer createCommandBuffer(const wgpu::TextureView backbufferView, size_t i) {
Kai Ninomiyac16a67a2017-07-27 18:30:57 -0700271 auto& bufferDst = particleBuffers[(i + 1) % 2];
Corentin Wallez04863c42019-10-25 11:36:47 +0000272 wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
Corentin Wallezf07e3bd2017-04-20 14:38:20 -0400273
Corentin Wallez82fbccb2018-09-21 00:24:37 +0000274 {
Corentin Wallez04863c42019-10-25 11:36:47 +0000275 wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
Yan, Shaobo300eec02018-12-21 10:40:26 +0000276 pass.SetPipeline(updatePipeline);
Corentin Wallez70c8c102019-10-09 16:08:42 +0000277 pass.SetBindGroup(0, updateBGs[i]);
Corentin Wallez3da19b82020-03-31 16:23:35 +0000278 pass.Dispatch(kNumParticles);
Corentin Wallez82fbccb2018-09-21 00:24:37 +0000279 pass.EndPass();
280 }
Corentin Wallezf07e3bd2017-04-20 14:38:20 -0400281
Corentin Wallez82fbccb2018-09-21 00:24:37 +0000282 {
Corentin Wallez604072b2019-11-12 18:30:11 +0000283 utils::ComboRenderPassDescriptor renderPass({backbufferView}, depthStencilView);
Corentin Wallez04863c42019-10-25 11:36:47 +0000284 wgpu::RenderPassEncoder pass = encoder.BeginRenderPass(&renderPass);
Yan, Shaobo300eec02018-12-21 10:40:26 +0000285 pass.SetPipeline(renderPipeline);
François Beaufort91b21422019-10-10 07:29:58 +0000286 pass.SetVertexBuffer(0, bufferDst);
287 pass.SetVertexBuffer(1, modelBuffer);
Corentin Wallez67b1ad72020-03-31 16:21:35 +0000288 pass.Draw(3, kNumParticles);
Corentin Wallez82fbccb2018-09-21 00:24:37 +0000289 pass.EndPass();
290 }
291
Corentin Walleze1f0d4e2019-02-15 12:54:08 +0000292 return encoder.Finish();
Corentin Wallezf07e3bd2017-04-20 14:38:20 -0400293}
294
295void init() {
Corentin Wallez39039fa2018-07-18 14:06:10 +0200296 device = CreateCppDawnDevice();
Corentin Wallezf07e3bd2017-04-20 14:38:20 -0400297
Corentin Wallez6d315da2021-02-04 15:33:42 +0000298 queue = device.GetQueue();
Kai Ninomiyac16a67a2017-07-27 18:30:57 -0700299 swapchain = GetSwapChain(device);
Corentin Wallez6b087812020-10-27 15:35:56 +0000300 swapchain.Configure(GetPreferredSwapChainTextureFormat(), wgpu::TextureUsage::RenderAttachment,
Corentin Wallez9e9e29f2019-08-27 08:21:39 +0000301 640, 480);
Corentin Wallezf07e3bd2017-04-20 14:38:20 -0400302
303 initBuffers();
304 initRender();
305 initSim();
Corentin Wallezf07e3bd2017-04-20 14:38:20 -0400306}
307
308void frame() {
Corentin Wallez604072b2019-11-12 18:30:11 +0000309 wgpu::TextureView backbufferView = swapchain.GetCurrentTextureView();
Kai Ninomiyac16a67a2017-07-27 18:30:57 -0700310
Corentin Wallez604072b2019-11-12 18:30:11 +0000311 wgpu::CommandBuffer commandBuffer = createCommandBuffer(backbufferView, pingpong);
Kai Ninomiyac16a67a2017-07-27 18:30:57 -0700312 queue.Submit(1, &commandBuffer);
Corentin Wallez604072b2019-11-12 18:30:11 +0000313 swapchain.Present();
Kai Ninomiyac16a67a2017-07-27 18:30:57 -0700314 DoFlush();
Corentin Wallezf07e3bd2017-04-20 14:38:20 -0400315
316 pingpong = (pingpong + 1) % 2;
317}
318
319int main(int argc, const char* argv[]) {
Corentin Wallez9347e8f2017-06-19 13:15:13 -0400320 if (!InitSample(argc, argv)) {
Corentin Wallezf07e3bd2017-04-20 14:38:20 -0400321 return 1;
322 }
323 init();
324
325 while (!ShouldQuit()) {
Austin Eng700a5fb2021-06-24 19:21:31 +0000326 utils::ScopedAutoreleasePool pool;
Corentin Wallezf07e3bd2017-04-20 14:38:20 -0400327 frame();
Corentin Wallez134e0802017-07-17 17:13:57 -0400328 utils::USleep(16000);
Corentin Wallezf07e3bd2017-04-20 14:38:20 -0400329 }
330
331 // TODO release stuff
332}