From 92881224ea03bf0088910f8822d8ecca40b6ae59 Mon Sep 17 00:00:00 2001 From: Tibo De Peuter Date: Fri, 29 Nov 2024 14:52:31 +0100 Subject: [PATCH 1/2] chore(1): Remove debugging print statements --- src/GpuPerlinNoise.cpp | 11 +++-------- src/PerlinNoise.cl | 2 -- 2 files changed, 3 insertions(+), 10 deletions(-) diff --git a/src/GpuPerlinNoise.cpp b/src/GpuPerlinNoise.cpp index 76a5999..7bccc32 100644 --- a/src/GpuPerlinNoise.cpp +++ b/src/GpuPerlinNoise.cpp @@ -119,14 +119,9 @@ void GpuPerlinNoise::Run() { float globalMin = *std::min_element(workgroupMin.begin(), workgroupMin.end()); float globalMax = *std::max_element(workgroupMax.begin(), workgroupMax.end()); /* Rescale outputNoise to [0, 255] */ - for (int i = 0; i < nrPixels; i++) { - outputImage[i] = static_cast(255 * (outputNoise[i] - globalMin) / (globalMax - globalMin)); - } - - // -------------- - printf("Debugging: min: %f, max: %f\n", workgroupMin[0], workgroupMax[0]); - printf("Debugging: gloabl_min: %f, global_max: %f\n", globalMin, globalMax); - // -------------- + for (int i = 0; i < nrPixels; i++) { + outputImage[i] = static_cast(255 * (outputNoise[i] - globalMin) / (globalMax - globalMin)); + } // write to PNG std::string pngPath = "perlin.png"; diff --git a/src/PerlinNoise.cl b/src/PerlinNoise.cl index 0c43a5d..ac93877 100644 --- a/src/PerlinNoise.cl +++ b/src/PerlinNoise.cl @@ -38,8 +38,6 @@ void kernel perlin_noise_texture(const int width, const int height, global float const float scale = 0.05f; // feel free to play with this float2 xy = (float2)(col * scale, row * scale); // (x,y) input for perlin noise algorithm - printf("col: %d, row: %d, pixeldId: %d, threadId: %d, workgroupId: %d\n", col, row, pixelId, threadId, workgroupId); - local float workgroup_p[16*16]; float x0 = floor(xy.x); From 31cc8f0a8457648576dc49cbbadec5ed9ba6acee Mon Sep 17 00:00:00 2001 From: Tibo De Peuter Date: Fri, 29 Nov 2024 14:53:04 +0100 Subject: [PATCH 2/2] 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