1cb93a386Sopenharmony_ci// Copyright 2017 The Dawn Authors
2cb93a386Sopenharmony_ci//
3cb93a386Sopenharmony_ci// Licensed under the Apache License, Version 2.0 (the "License");
4cb93a386Sopenharmony_ci// you may not use this file except in compliance with the License.
5cb93a386Sopenharmony_ci// You may obtain a copy of the License at
6cb93a386Sopenharmony_ci//
7cb93a386Sopenharmony_ci//     http://www.apache.org/licenses/LICENSE-2.0
8cb93a386Sopenharmony_ci//
9cb93a386Sopenharmony_ci// Unless required by applicable law or agreed to in writing, software
10cb93a386Sopenharmony_ci// distributed under the License is distributed on an "AS IS" BASIS,
11cb93a386Sopenharmony_ci// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
12cb93a386Sopenharmony_ci// See the License for the specific language governing permissions and
13cb93a386Sopenharmony_ci// limitations under the License.
14cb93a386Sopenharmony_ci
15cb93a386Sopenharmony_ci#include "SampleUtils.h"
16cb93a386Sopenharmony_ci
17cb93a386Sopenharmony_ci#include "utils/ComboRenderPipelineDescriptor.h"
18cb93a386Sopenharmony_ci#include "utils/ScopedAutoreleasePool.h"
19cb93a386Sopenharmony_ci#include "utils/SystemUtils.h"
20cb93a386Sopenharmony_ci#include "utils/WGPUHelpers.h"
21cb93a386Sopenharmony_ci
22cb93a386Sopenharmony_ci#include <array>
23cb93a386Sopenharmony_ci#include <cstring>
24cb93a386Sopenharmony_ci#include <random>
25cb93a386Sopenharmony_ci
26cb93a386Sopenharmony_ci#include <glm/glm.hpp>
27cb93a386Sopenharmony_ci
28cb93a386Sopenharmony_ciwgpu::Device device;
29cb93a386Sopenharmony_ciwgpu::Queue queue;
30cb93a386Sopenharmony_ciwgpu::SwapChain swapchain;
31cb93a386Sopenharmony_ciwgpu::TextureView depthStencilView;
32cb93a386Sopenharmony_ci
33cb93a386Sopenharmony_ciwgpu::Buffer modelBuffer;
34cb93a386Sopenharmony_cistd::array<wgpu::Buffer, 2> particleBuffers;
35cb93a386Sopenharmony_ci
36cb93a386Sopenharmony_ciwgpu::RenderPipeline renderPipeline;
37cb93a386Sopenharmony_ci
38cb93a386Sopenharmony_ciwgpu::Buffer updateParams;
39cb93a386Sopenharmony_ciwgpu::ComputePipeline updatePipeline;
40cb93a386Sopenharmony_cistd::array<wgpu::BindGroup, 2> updateBGs;
41cb93a386Sopenharmony_ci
42cb93a386Sopenharmony_cisize_t pingpong = 0;
43cb93a386Sopenharmony_ci
44cb93a386Sopenharmony_cistatic const uint32_t kNumParticles = 1000;
45cb93a386Sopenharmony_ci
46cb93a386Sopenharmony_cistruct Particle {
47cb93a386Sopenharmony_ci    glm::vec2 pos;
48cb93a386Sopenharmony_ci    glm::vec2 vel;
49cb93a386Sopenharmony_ci};
50cb93a386Sopenharmony_ci
51cb93a386Sopenharmony_cistruct SimParams {
52cb93a386Sopenharmony_ci    float deltaT;
53cb93a386Sopenharmony_ci    float rule1Distance;
54cb93a386Sopenharmony_ci    float rule2Distance;
55cb93a386Sopenharmony_ci    float rule3Distance;
56cb93a386Sopenharmony_ci    float rule1Scale;
57cb93a386Sopenharmony_ci    float rule2Scale;
58cb93a386Sopenharmony_ci    float rule3Scale;
59cb93a386Sopenharmony_ci    int particleCount;
60cb93a386Sopenharmony_ci};
61cb93a386Sopenharmony_ci
62cb93a386Sopenharmony_civoid initBuffers() {
63cb93a386Sopenharmony_ci    glm::vec2 model[3] = {
64cb93a386Sopenharmony_ci        {-0.01, -0.02},
65cb93a386Sopenharmony_ci        {0.01, -0.02},
66cb93a386Sopenharmony_ci        {0.00, 0.02},
67cb93a386Sopenharmony_ci    };
68cb93a386Sopenharmony_ci    modelBuffer =
69cb93a386Sopenharmony_ci        utils::CreateBufferFromData(device, model, sizeof(model), wgpu::BufferUsage::Vertex);
70cb93a386Sopenharmony_ci
71cb93a386Sopenharmony_ci    SimParams params = {0.04f, 0.1f, 0.025f, 0.025f, 0.02f, 0.05f, 0.005f, kNumParticles};
72cb93a386Sopenharmony_ci    updateParams =
73cb93a386Sopenharmony_ci        utils::CreateBufferFromData(device, &params, sizeof(params), wgpu::BufferUsage::Uniform);
74cb93a386Sopenharmony_ci
75cb93a386Sopenharmony_ci    std::vector<Particle> initialParticles(kNumParticles);
76cb93a386Sopenharmony_ci    {
77cb93a386Sopenharmony_ci        std::mt19937 generator;
78cb93a386Sopenharmony_ci        std::uniform_real_distribution<float> dist(-1.0f, 1.0f);
79cb93a386Sopenharmony_ci        for (auto& p : initialParticles) {
80cb93a386Sopenharmony_ci            p.pos = glm::vec2(dist(generator), dist(generator));
81cb93a386Sopenharmony_ci            p.vel = glm::vec2(dist(generator), dist(generator)) * 0.1f;
82cb93a386Sopenharmony_ci        }
83cb93a386Sopenharmony_ci    }
84cb93a386Sopenharmony_ci
85cb93a386Sopenharmony_ci    for (size_t i = 0; i < 2; i++) {
86cb93a386Sopenharmony_ci        wgpu::BufferDescriptor descriptor;
87cb93a386Sopenharmony_ci        descriptor.size = sizeof(Particle) * kNumParticles;
88cb93a386Sopenharmony_ci        descriptor.usage =
89cb93a386Sopenharmony_ci            wgpu::BufferUsage::CopyDst | wgpu::BufferUsage::Vertex | wgpu::BufferUsage::Storage;
90cb93a386Sopenharmony_ci        particleBuffers[i] = device.CreateBuffer(&descriptor);
91cb93a386Sopenharmony_ci
92cb93a386Sopenharmony_ci        queue.WriteBuffer(particleBuffers[i], 0,
93cb93a386Sopenharmony_ci                          reinterpret_cast<uint8_t*>(initialParticles.data()),
94cb93a386Sopenharmony_ci                          sizeof(Particle) * kNumParticles);
95cb93a386Sopenharmony_ci    }
96cb93a386Sopenharmony_ci}
97cb93a386Sopenharmony_ci
98cb93a386Sopenharmony_civoid initRender() {
99cb93a386Sopenharmony_ci    wgpu::ShaderModule vsModule = utils::CreateShaderModule(device, R"(
100cb93a386Sopenharmony_ci        struct VertexIn {
101cb93a386Sopenharmony_ci            [[location(0)]] a_particlePos : vec2<f32>;
102cb93a386Sopenharmony_ci            [[location(1)]] a_particleVel : vec2<f32>;
103cb93a386Sopenharmony_ci            [[location(2)]] a_pos : vec2<f32>;
104cb93a386Sopenharmony_ci        };
105cb93a386Sopenharmony_ci
106cb93a386Sopenharmony_ci        [[stage(vertex)]]
107cb93a386Sopenharmony_ci        fn main(input : VertexIn) -> [[builtin(position)]] vec4<f32> {
108cb93a386Sopenharmony_ci            var angle : f32 = -atan2(input.a_particleVel.x, input.a_particleVel.y);
109cb93a386Sopenharmony_ci            var pos : vec2<f32> = vec2<f32>(
110cb93a386Sopenharmony_ci                (input.a_pos.x * cos(angle)) - (input.a_pos.y * sin(angle)),
111cb93a386Sopenharmony_ci                (input.a_pos.x * sin(angle)) + (input.a_pos.y * cos(angle)));
112cb93a386Sopenharmony_ci            return vec4<f32>(pos + input.a_particlePos, 0.0, 1.0);
113cb93a386Sopenharmony_ci        }
114cb93a386Sopenharmony_ci    )");
115cb93a386Sopenharmony_ci
116cb93a386Sopenharmony_ci    wgpu::ShaderModule fsModule = utils::CreateShaderModule(device, R"(
117cb93a386Sopenharmony_ci        [[stage(fragment)]]
118cb93a386Sopenharmony_ci        fn main() -> [[location(0)]] vec4<f32> {
119cb93a386Sopenharmony_ci            return vec4<f32>(1.0, 1.0, 1.0, 1.0);
120cb93a386Sopenharmony_ci        }
121cb93a386Sopenharmony_ci    )");
122cb93a386Sopenharmony_ci
123cb93a386Sopenharmony_ci    depthStencilView = CreateDefaultDepthStencilView(device);
124cb93a386Sopenharmony_ci
125cb93a386Sopenharmony_ci    utils::ComboRenderPipelineDescriptor descriptor;
126cb93a386Sopenharmony_ci
127cb93a386Sopenharmony_ci    descriptor.vertex.module = vsModule;
128cb93a386Sopenharmony_ci    descriptor.vertex.bufferCount = 2;
129cb93a386Sopenharmony_ci    descriptor.cBuffers[0].arrayStride = sizeof(Particle);
130cb93a386Sopenharmony_ci    descriptor.cBuffers[0].stepMode = wgpu::VertexStepMode::Instance;
131cb93a386Sopenharmony_ci    descriptor.cBuffers[0].attributeCount = 2;
132cb93a386Sopenharmony_ci    descriptor.cAttributes[0].offset = offsetof(Particle, pos);
133cb93a386Sopenharmony_ci    descriptor.cAttributes[0].format = wgpu::VertexFormat::Float32x2;
134cb93a386Sopenharmony_ci    descriptor.cAttributes[1].shaderLocation = 1;
135cb93a386Sopenharmony_ci    descriptor.cAttributes[1].offset = offsetof(Particle, vel);
136cb93a386Sopenharmony_ci    descriptor.cAttributes[1].format = wgpu::VertexFormat::Float32x2;
137cb93a386Sopenharmony_ci    descriptor.cBuffers[1].arrayStride = sizeof(glm::vec2);
138cb93a386Sopenharmony_ci    descriptor.cBuffers[1].attributeCount = 1;
139cb93a386Sopenharmony_ci    descriptor.cBuffers[1].attributes = &descriptor.cAttributes[2];
140cb93a386Sopenharmony_ci    descriptor.cAttributes[2].shaderLocation = 2;
141cb93a386Sopenharmony_ci    descriptor.cAttributes[2].format = wgpu::VertexFormat::Float32x2;
142cb93a386Sopenharmony_ci
143cb93a386Sopenharmony_ci    descriptor.cFragment.module = fsModule;
144cb93a386Sopenharmony_ci    descriptor.EnableDepthStencil(wgpu::TextureFormat::Depth24PlusStencil8);
145cb93a386Sopenharmony_ci    descriptor.cTargets[0].format = GetPreferredSwapChainTextureFormat();
146cb93a386Sopenharmony_ci
147cb93a386Sopenharmony_ci    renderPipeline = device.CreateRenderPipeline(&descriptor);
148cb93a386Sopenharmony_ci}
149cb93a386Sopenharmony_ci
150cb93a386Sopenharmony_civoid initSim() {
151cb93a386Sopenharmony_ci    wgpu::ShaderModule module = utils::CreateShaderModule(device, R"(
152cb93a386Sopenharmony_ci        struct Particle {
153cb93a386Sopenharmony_ci            pos : vec2<f32>;
154cb93a386Sopenharmony_ci            vel : vec2<f32>;
155cb93a386Sopenharmony_ci        };
156cb93a386Sopenharmony_ci        [[block]] struct SimParams {
157cb93a386Sopenharmony_ci            deltaT : f32;
158cb93a386Sopenharmony_ci            rule1Distance : f32;
159cb93a386Sopenharmony_ci            rule2Distance : f32;
160cb93a386Sopenharmony_ci            rule3Distance : f32;
161cb93a386Sopenharmony_ci            rule1Scale : f32;
162cb93a386Sopenharmony_ci            rule2Scale : f32;
163cb93a386Sopenharmony_ci            rule3Scale : f32;
164cb93a386Sopenharmony_ci            particleCount : u32;
165cb93a386Sopenharmony_ci        };
166cb93a386Sopenharmony_ci        [[block]] struct Particles {
167cb93a386Sopenharmony_ci            particles : array<Particle>;
168cb93a386Sopenharmony_ci        };
169cb93a386Sopenharmony_ci        [[binding(0), group(0)]] var<uniform> params : SimParams;
170cb93a386Sopenharmony_ci        [[binding(1), group(0)]] var<storage, read> particlesA : Particles;
171cb93a386Sopenharmony_ci        [[binding(2), group(0)]] var<storage, read_write> particlesB : Particles;
172cb93a386Sopenharmony_ci
173cb93a386Sopenharmony_ci        // https://github.com/austinEng/Project6-Vulkan-Flocking/blob/master/data/shaders/computeparticles/particle.comp
174cb93a386Sopenharmony_ci        [[stage(compute), workgroup_size(1)]]
175cb93a386Sopenharmony_ci        fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
176cb93a386Sopenharmony_ci            var index : u32 = GlobalInvocationID.x;
177cb93a386Sopenharmony_ci            if (index >= params.particleCount) {
178cb93a386Sopenharmony_ci                return;
179cb93a386Sopenharmony_ci            }
180cb93a386Sopenharmony_ci            var vPos : vec2<f32> = particlesA.particles[index].pos;
181cb93a386Sopenharmony_ci            var vVel : vec2<f32> = particlesA.particles[index].vel;
182cb93a386Sopenharmony_ci            var cMass : vec2<f32> = vec2<f32>(0.0, 0.0);
183cb93a386Sopenharmony_ci            var cVel : vec2<f32> = vec2<f32>(0.0, 0.0);
184cb93a386Sopenharmony_ci            var colVel : vec2<f32> = vec2<f32>(0.0, 0.0);
185cb93a386Sopenharmony_ci            var cMassCount : u32 = 0u;
186cb93a386Sopenharmony_ci            var cVelCount : u32 = 0u;
187cb93a386Sopenharmony_ci            var pos : vec2<f32>;
188cb93a386Sopenharmony_ci            var vel : vec2<f32>;
189cb93a386Sopenharmony_ci
190cb93a386Sopenharmony_ci            for (var i : u32 = 0u; i < params.particleCount; i = i + 1u) {
191cb93a386Sopenharmony_ci                if (i == index) {
192cb93a386Sopenharmony_ci                    continue;
193cb93a386Sopenharmony_ci                }
194cb93a386Sopenharmony_ci
195cb93a386Sopenharmony_ci                pos = particlesA.particles[i].pos.xy;
196cb93a386Sopenharmony_ci                vel = particlesA.particles[i].vel.xy;
197cb93a386Sopenharmony_ci                if (distance(pos, vPos) < params.rule1Distance) {
198cb93a386Sopenharmony_ci                    cMass = cMass + pos;
199cb93a386Sopenharmony_ci                    cMassCount = cMassCount + 1u;
200cb93a386Sopenharmony_ci                }
201cb93a386Sopenharmony_ci                if (distance(pos, vPos) < params.rule2Distance) {
202cb93a386Sopenharmony_ci                    colVel = colVel - (pos - vPos);
203cb93a386Sopenharmony_ci                }
204cb93a386Sopenharmony_ci                if (distance(pos, vPos) < params.rule3Distance) {
205cb93a386Sopenharmony_ci                    cVel = cVel + vel;
206cb93a386Sopenharmony_ci                    cVelCount = cVelCount + 1u;
207cb93a386Sopenharmony_ci                }
208cb93a386Sopenharmony_ci            }
209cb93a386Sopenharmony_ci
210cb93a386Sopenharmony_ci            if (cMassCount > 0u) {
211cb93a386Sopenharmony_ci                cMass = (cMass / vec2<f32>(f32(cMassCount), f32(cMassCount))) - vPos;
212cb93a386Sopenharmony_ci            }
213cb93a386Sopenharmony_ci
214cb93a386Sopenharmony_ci            if (cVelCount > 0u) {
215cb93a386Sopenharmony_ci                cVel = cVel / vec2<f32>(f32(cVelCount), f32(cVelCount));
216cb93a386Sopenharmony_ci            }
217cb93a386Sopenharmony_ci            vVel = vVel + (cMass * params.rule1Scale) + (colVel * params.rule2Scale) +
218cb93a386Sopenharmony_ci                (cVel * params.rule3Scale);
219cb93a386Sopenharmony_ci
220cb93a386Sopenharmony_ci            // clamp velocity for a more pleasing simulation
221cb93a386Sopenharmony_ci            vVel = normalize(vVel) * clamp(length(vVel), 0.0, 0.1);
222cb93a386Sopenharmony_ci            // kinematic update
223cb93a386Sopenharmony_ci            vPos = vPos + (vVel * params.deltaT);
224cb93a386Sopenharmony_ci
225cb93a386Sopenharmony_ci            // Wrap around boundary
226cb93a386Sopenharmony_ci            if (vPos.x < -1.0) {
227cb93a386Sopenharmony_ci                vPos.x = 1.0;
228cb93a386Sopenharmony_ci            }
229cb93a386Sopenharmony_ci            if (vPos.x > 1.0) {
230cb93a386Sopenharmony_ci                vPos.x = -1.0;
231cb93a386Sopenharmony_ci            }
232cb93a386Sopenharmony_ci            if (vPos.y < -1.0) {
233cb93a386Sopenharmony_ci                vPos.y = 1.0;
234cb93a386Sopenharmony_ci            }
235cb93a386Sopenharmony_ci            if (vPos.y > 1.0) {
236cb93a386Sopenharmony_ci                vPos.y = -1.0;
237cb93a386Sopenharmony_ci            }
238cb93a386Sopenharmony_ci
239cb93a386Sopenharmony_ci            // Write back
240cb93a386Sopenharmony_ci            particlesB.particles[index].pos = vPos;
241cb93a386Sopenharmony_ci            particlesB.particles[index].vel = vVel;
242cb93a386Sopenharmony_ci            return;
243cb93a386Sopenharmony_ci        }
244cb93a386Sopenharmony_ci    )");
245cb93a386Sopenharmony_ci
246cb93a386Sopenharmony_ci    auto bgl = utils::MakeBindGroupLayout(
247cb93a386Sopenharmony_ci        device, {
248cb93a386Sopenharmony_ci                    {0, wgpu::ShaderStage::Compute, wgpu::BufferBindingType::Uniform},
249cb93a386Sopenharmony_ci                    {1, wgpu::ShaderStage::Compute, wgpu::BufferBindingType::Storage},
250cb93a386Sopenharmony_ci                    {2, wgpu::ShaderStage::Compute, wgpu::BufferBindingType::Storage},
251cb93a386Sopenharmony_ci                });
252cb93a386Sopenharmony_ci
253cb93a386Sopenharmony_ci    wgpu::PipelineLayout pl = utils::MakeBasicPipelineLayout(device, &bgl);
254cb93a386Sopenharmony_ci
255cb93a386Sopenharmony_ci    wgpu::ComputePipelineDescriptor csDesc;
256cb93a386Sopenharmony_ci    csDesc.layout = pl;
257cb93a386Sopenharmony_ci    csDesc.compute.module = module;
258cb93a386Sopenharmony_ci    csDesc.compute.entryPoint = "main";
259cb93a386Sopenharmony_ci    updatePipeline = device.CreateComputePipeline(&csDesc);
260cb93a386Sopenharmony_ci
261cb93a386Sopenharmony_ci    for (uint32_t i = 0; i < 2; ++i) {
262cb93a386Sopenharmony_ci        updateBGs[i] = utils::MakeBindGroup(
263cb93a386Sopenharmony_ci            device, bgl,
264cb93a386Sopenharmony_ci            {
265cb93a386Sopenharmony_ci                {0, updateParams, 0, sizeof(SimParams)},
266cb93a386Sopenharmony_ci                {1, particleBuffers[i], 0, kNumParticles * sizeof(Particle)},
267cb93a386Sopenharmony_ci                {2, particleBuffers[(i + 1) % 2], 0, kNumParticles * sizeof(Particle)},
268cb93a386Sopenharmony_ci            });
269cb93a386Sopenharmony_ci    }
270cb93a386Sopenharmony_ci}
271cb93a386Sopenharmony_ci
272cb93a386Sopenharmony_ciwgpu::CommandBuffer createCommandBuffer(const wgpu::TextureView backbufferView, size_t i) {
273cb93a386Sopenharmony_ci    auto& bufferDst = particleBuffers[(i + 1) % 2];
274cb93a386Sopenharmony_ci    wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
275cb93a386Sopenharmony_ci
276cb93a386Sopenharmony_ci    {
277cb93a386Sopenharmony_ci        wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
278cb93a386Sopenharmony_ci        pass.SetPipeline(updatePipeline);
279cb93a386Sopenharmony_ci        pass.SetBindGroup(0, updateBGs[i]);
280cb93a386Sopenharmony_ci        pass.Dispatch(kNumParticles);
281cb93a386Sopenharmony_ci        pass.EndPass();
282cb93a386Sopenharmony_ci    }
283cb93a386Sopenharmony_ci
284cb93a386Sopenharmony_ci    {
285cb93a386Sopenharmony_ci        utils::ComboRenderPassDescriptor renderPass({backbufferView}, depthStencilView);
286cb93a386Sopenharmony_ci        wgpu::RenderPassEncoder pass = encoder.BeginRenderPass(&renderPass);
287cb93a386Sopenharmony_ci        pass.SetPipeline(renderPipeline);
288cb93a386Sopenharmony_ci        pass.SetVertexBuffer(0, bufferDst);
289cb93a386Sopenharmony_ci        pass.SetVertexBuffer(1, modelBuffer);
290cb93a386Sopenharmony_ci        pass.Draw(3, kNumParticles);
291cb93a386Sopenharmony_ci        pass.EndPass();
292cb93a386Sopenharmony_ci    }
293cb93a386Sopenharmony_ci
294cb93a386Sopenharmony_ci    return encoder.Finish();
295cb93a386Sopenharmony_ci}
296cb93a386Sopenharmony_ci
297cb93a386Sopenharmony_civoid init() {
298cb93a386Sopenharmony_ci    device = CreateCppDawnDevice();
299cb93a386Sopenharmony_ci
300cb93a386Sopenharmony_ci    queue = device.GetQueue();
301cb93a386Sopenharmony_ci    swapchain = GetSwapChain(device);
302cb93a386Sopenharmony_ci    swapchain.Configure(GetPreferredSwapChainTextureFormat(), wgpu::TextureUsage::RenderAttachment,
303cb93a386Sopenharmony_ci                        640, 480);
304cb93a386Sopenharmony_ci
305cb93a386Sopenharmony_ci    initBuffers();
306cb93a386Sopenharmony_ci    initRender();
307cb93a386Sopenharmony_ci    initSim();
308cb93a386Sopenharmony_ci}
309cb93a386Sopenharmony_ci
310cb93a386Sopenharmony_civoid frame() {
311cb93a386Sopenharmony_ci    wgpu::TextureView backbufferView = swapchain.GetCurrentTextureView();
312cb93a386Sopenharmony_ci
313cb93a386Sopenharmony_ci    wgpu::CommandBuffer commandBuffer = createCommandBuffer(backbufferView, pingpong);
314cb93a386Sopenharmony_ci    queue.Submit(1, &commandBuffer);
315cb93a386Sopenharmony_ci    swapchain.Present();
316cb93a386Sopenharmony_ci    DoFlush();
317cb93a386Sopenharmony_ci
318cb93a386Sopenharmony_ci    pingpong = (pingpong + 1) % 2;
319cb93a386Sopenharmony_ci}
320cb93a386Sopenharmony_ci
321cb93a386Sopenharmony_ciint main(int argc, const char* argv[]) {
322cb93a386Sopenharmony_ci    if (!InitSample(argc, argv)) {
323cb93a386Sopenharmony_ci        return 1;
324cb93a386Sopenharmony_ci    }
325cb93a386Sopenharmony_ci    init();
326cb93a386Sopenharmony_ci
327cb93a386Sopenharmony_ci    while (!ShouldQuit()) {
328cb93a386Sopenharmony_ci        utils::ScopedAutoreleasePool pool;
329cb93a386Sopenharmony_ci        frame();
330cb93a386Sopenharmony_ci        utils::USleep(16000);
331cb93a386Sopenharmony_ci    }
332cb93a386Sopenharmony_ci
333cb93a386Sopenharmony_ci    // TODO release stuff
334cb93a386Sopenharmony_ci}
335