From 31cc8f0a8457648576dc49cbbadec5ed9ba6acee Mon Sep 17 00:00:00 2001 From: Tibo De Peuter Date: Fri, 29 Nov 2024 14:53:04 +0100 Subject: [PATCH] feat(2): Assignment 2 --- src/GpuParticleSystem.cpp | 33 ++++++++++++++ src/ParticleSystem.cl | 91 ++++++++++++++++++++++++++++++++++++++- 2 files changed, 123 insertions(+), 1 deletion(-) diff --git a/src/GpuParticleSystem.cpp b/src/GpuParticleSystem.cpp index 0bd2fd7..7d50c61 100644 --- a/src/GpuParticleSystem.cpp +++ b/src/GpuParticleSystem.cpp @@ -1,6 +1,9 @@ #include "GpuApplication.h" +#define BOUNDARY 0.9f +#define PUSBACK 0.05f + GpuParticleSystem::GpuParticleSystem(std::string kernelFile) { this->kernelFile = kernelFile; srand(0); @@ -78,6 +81,10 @@ void GpuParticleSystem::InitPosVel() { // -------------- // Your code here // Initialize positions and velocities so that the particles start off on the left side of the screen. + for (int i = 0; i < nrParticles; i++) { + positions[i * 2] = -BOUNDARY + PUSBACK; + positions[i * 2 + 1] = -BOUNDARY + PUSBACK + (2 * BOUNDARY - 2 * PUSBACK) * i / nrParticles; + } // -------------- } @@ -88,6 +95,16 @@ void GpuParticleSystem::Run() { // -------------- // Your code here // Setup OpenCL so that the 'updateParticles' kernel can be executed for every frame in the loop below + cl::Buffer positions_device(context, CL_MEM_READ_WRITE, sizeof(float) * nrParticles * 2); + cl::Buffer velocities_device(context, CL_MEM_READ_WRITE, sizeof(float) * nrParticles * 2); + + cl::KernelFunctor kernelFunctor(program, "updateParticles"); + cl::NDRange rangeGlobal(nrParticles); + cl::EnqueueArgs enqueArgs(queue, rangeGlobal); + + // Write initial data to the device buffers. + queue.enqueueWriteBuffer(positions_device, CL_TRUE, 0, sizeof(float) * nrParticles * 2, positions.data()); + queue.enqueueWriteBuffer(velocities_device, CL_TRUE, 0, sizeof(float) * nrParticles * 2, velocities.data()); // -------------- // Setup the OpenGL VBOs that will pass the position and velocity of each particle to the vertex shader @@ -123,6 +140,9 @@ void GpuParticleSystem::Run() { // -------------- // Your code here // Update the OpenCL cl::Buffers holding the positions and velocities + printf("Resetting the simulation\n"); + queue.enqueueWriteBuffer(positions_device, CL_TRUE, 0, sizeof(float) * nrParticles * 2, positions.data()); + queue.enqueueWriteBuffer(velocities_device, CL_TRUE, 0, sizeof(float) * nrParticles * 2, velocities.data()); // -------------- reset = false; } @@ -131,6 +151,19 @@ void GpuParticleSystem::Run() { // Your code here // Update the particles' position and velocities using the OpenCL kernel // Make sure that the OpenGL VBOs are updated with these new values (hint: glBindBuffer, glBufferSubData ) + + // Launch the kernel + cl::Event event = kernelFunctor(enqueArgs, positions_device, velocities_device, dt, cursorX, cursorY); + + // Read the data from the device buffer into the host vector + queue.enqueueReadBuffer(positions_device, CL_TRUE, 0, sizeof(float) * nrParticles * 2, positions.data()); + queue.enqueueReadBuffer(velocities_device, CL_TRUE, 0, sizeof(float) * nrParticles * 2, velocities.data()); + + // Update the OpenGL VBOs + glBindBuffer(GL_ARRAY_BUFFER, VBOs[0]); + glBufferSubData(GL_ARRAY_BUFFER, 0, sizeof(float) * nrParticles * 2, positions.data()); + glBindBuffer(GL_ARRAY_BUFFER, VBOs[1]); + glBufferSubData(GL_ARRAY_BUFFER, 0, sizeof(float) * nrParticles * 2, velocities.data()); // -------------- // render to screen diff --git a/src/ParticleSystem.cl b/src/ParticleSystem.cl index 94182bf..a73baf6 100644 --- a/src/ParticleSystem.cl +++ b/src/ParticleSystem.cl @@ -5,7 +5,7 @@ // ------------------------------------- void kernel updateParticles(global float* positions, global float* velocities, float dt, float cursorX, float cursorY) { - + // feel free to play with these values float particleRepulsionRadius = 0.1f; float particleRepulsionStrength = 10.0f; @@ -14,4 +14,93 @@ void kernel updateParticles(global float* positions, global float* velocities, f float damping = 0.96f; // Your code here + const float boundary = 0.9f; + const float pushback = 0.05f; + const float boundaryScale = 1 / (boundaryRepulsionRadius * boundaryRepulsionRadius) * boundaryRepulsionStrength; + const float particleScale = 1/ particleRepulsionRadius * particleRepulsionStrength; + + /* 1. Get the index of the current thread, and thus the indices of the current particle in + positions and velocities arrays. */ + const int particleId = get_global_id(0); + const int x = particleId * 2; + const int y = particleId * 2 + 1; + + /* 4. The particle should be affected by gravity. Use g = 9.8 m/s^2. */ + velocities[y] -= 9.8f * dt; + + /* 5. Several repulsion forces should work on the particle. Note however, that force = m * a, + but we neglect mass here, so instead of applying N forces on the particle, we apply N + acceleration vectors, which, when all summed up change the velocities (using dt). */ + /* 5.1. The particle should stay within the boundary box. */ + // When a particle crosses a boundary, it should be teleported back to the correct side of the + // boundary, and the velocity along this axis should be reset to 0. + if (positions[x] < -boundary) { + positions[x] = -boundary + pushback; + velocities[x] = 0; + } else if (positions[x] > boundary) { + positions[x] = boundary - pushback; + velocities[x] = 0; + } + if (positions[y] < -boundary) { + positions[y] = -boundary + pushback; + velocities[y] = 0; + } else if (positions[y] > boundary) { + positions[y] = boundary - pushback; + velocities[y] = 0; + } + // Secondly, if a particle is within boundaryRepulsionRadius of a boundary, a repulsion force + // should be added, pointed away from the boundary. The acceleration should increase according + // to a parabole. (f(x) = x^2) + if (positions[x] < -boundary + boundaryRepulsionRadius) { + float pos = -boundary + boundaryRepulsionRadius - positions[x]; + velocities[x] += boundaryScale * pos * pos * dt; + } else if (positions[x] > boundary - boundaryRepulsionRadius) { + float pos = positions[x] - boundary + boundaryRepulsionRadius; + velocities[x] -= boundaryScale * pos * pos * dt; + } + if (positions[y] < -boundary + boundaryRepulsionRadius) { + float pos = -boundary + boundaryRepulsionRadius - positions[y]; + velocities[y] += boundaryScale * pos * pos * dt; + } else if (positions[y] > boundary - boundaryRepulsionRadius) { + float pos = positions[y] - boundary + boundaryRepulsionRadius; + velocities[y] -= boundaryScale * pos * pos * dt; + } + /* 5.2. The particle should be repelled from other particles within repulsionRadius. Similarly + to the previous point, you can use repulsionRadius and repulsionStrength. Instead of a + parabole, you can simply make the norm of the acceleration vector increase linearly + with closeness to the other particle. */ + for (int i = 0; i < get_global_size(0); i++) { + if (i == particleId) { + continue; + } + float dx = positions[x] - positions[i * 2]; + float dy = positions[y] - positions[i * 2 + 1]; + float dist = sqrt(dx * dx + dy * dy); + if (dist < particleRepulsionRadius) { + float scale = particleScale * (particleRepulsionRadius - dist) / dist; + velocities[x] += scale * dx * dt; + velocities[y] += scale * dy * dt; + } + } + /* 5.3. The particle should be repelled from the cursor. You can reuse boundaryRepulsionRadius + and boundaryRepulsionStrength. */ + float dx = positions[x] - cursorX; + float dy = positions[y] - cursorY; + float dist = sqrt(dx * dx + dy * dy); + if (dist < boundaryRepulsionRadius) { + float scale = boundaryScale * (boundaryRepulsionRadius - dist) / dist; + velocities[x] += scale * dx * dt; + velocities[y] += scale * dy * dt; + } + + /* 3. Just before this, the velocity should be dampened as follows: */ + velocities[x] *= damping; + velocities[y] *= damping; + + /* 2. At the end of this function, the position of the particle should be updated using its + velocity and dt (= the duration of 1 frame, in seconds). Hopefully you still remember how + to update a position based on the velocity, and the velocity based on the acceleration, + and a small timestep (dt). */ + positions[x] += velocities[x] * dt; + positions[y] += velocities[y] * dt; } \ No newline at end of file