diff --git a/README.md b/README.md index 98dd9a8..f23dbe0 100644 --- a/README.md +++ b/README.md @@ -1,10 +1,65 @@ **University of Pennsylvania, CIS 565: GPU Programming and Architecture, Project 1 - Flocking** -* (TODO) YOUR NAME HERE -* Tested on: (TODO) Windows 22, i7-2222 @ 2.22GHz 22GB, GTX 222 222MB (Moore 2222 Lab) +* Joseph Klinger +* Tested on: Windows 10, i5-7300HQ (4 CPUs) @ ~2.50GHz 2B, GTX 1050 6030MB (Personal Machine) -### (TODO: Your README) +### README -Include screenshots, analysis, etc. (Remember, this is public, so don't put -anything here that you don't want to share with the world.) +![](boids.gif) + +In this project, we are attempting to simulate the motion of "Boids," which move each frame according to the properties of their neighbors. This can be done a number of ways, +one of the better ones being to utilize a spatial uniform grid to accelerate the processing of each boid being simulated. On the CPU (even with an acceleration structure such as a uniform grid in place), +we would be forced to iterate over each boid to process them individually, which is valid, but doesn't scale up well. So, we harness the power of the GPU through CUDA to simulate these Boids in parallel. + +The naive way of computing how a Boid should move is to check every other Boid in the simulation. However, Boids have a finite (and usually small) neighborhood radius, and only are affected by Boids within that radius. +Ideally, we do not have to bother checking faraway Boids that will have no effect on the Boid in question. As mentioned before, a Uniform Grid is utilized to narrow down the Boids that must be checked to only those in the +cells next to a given Boid. To further accelerate things, steps are taken to ensure that the arrays containing the Boid information (position and velocity) are arranged in the same way as the uniform grid indices for faster +read times - see INSTRUCTION.md for more. + +### Analysis + +**Data** + +I'm just going to throw some stats at you to start. Here are how the various implementations performed: + +![](graph1.png) + +Here we see that the naive implementation falls off rather quickly, as we would expect. The uniform grid makes the most significant optimization (again as we would expect), as we are cutting out +the vast majority of unnecessary-to-check Boids in the simulation. The coherent memory optimization, which allows the simulation to reduce the number of pointers dereferenced to global memory, is still appreciably better than +that, and moreso than I was expecting. The increased performance is welcome, though. + +![](graph2.png) + +Turning off the visualization allows for awesomely high performance, but the overall pattern remains the same. This jump in performance was actually much higher than I was anticipating - I knew draw calls in OpenGL could be +costly, but we can see that it really cut down on the overall performance of the simulation in the first graph. + +![](graph3.png) + +Here I analyzed the performance versus the number of threads per block for each method at 5000 Boids. As we will discuss below, the warp size on my laptop's GPU, the GTX 1050, is 32, so the performance plateaus after +increasing the blocksize to 32, as expected. + +**Questions:** +* Changing the number of boids: +Changing the number of boids, or the input "N" to our simulation, has the most direct effect on the performance of the simulation. +Each method (Naive, Uniform Grid, Coherent Grid) processes each boid to compute a new velocity (for that boid). +Each method must take into account a subset of the other boids to do this. So, the method that can take into account as few of +the other boids for a given boid will perform the best. Obviously, taking into account every other boid (Naive) will perform far worse +than taking into account only the nearby boids (Uniform Grid). + +* Changing the number of threads per block: +By changing the block size (and therefore block count), we are directly altering exactly how many boids are processed in parallel at a time. Given that on a GTX 1050, a maximum of 32 threads can be processed at once +in parallel, it makes sense that performance increase would plateau after increasing the number of threads per block to 32, as seen in the third graph above. Any number lower than 32, and the SMs aren't being used to their full +potential (not every thread in a warp of 32 threads would be active), and so we witness less performance. + +* Experiencing performance improvements with the more coherent uniform grid: +Comparing to the scattered uniform grid implemenation, adding memory coherence to the boid position/velocity arrays added a 10 - 40% performance increase (!!!) in certain cases, which is very appreciable! This was higher than +I was initially expecting, but it puts into perspective how costly it can be to derefence pointers to global memory repeatedly. In future projects, I will be sure to incorporate optimizations such as this. + +* Did changing cell width and checking 27 vs 8 neighboring cells: +Intuitively, we know that checking more cells than is necessary will negatively affect performance. Given that the grid cell size is two times the neighborhood distance, in three dimensions, only 8 grid cells even need to be checked. +However, here is a brief analysis on how the performance compared, for the coherent grid on a selected number of Boids: + +![](graph4.png) + +As expected, we experience a pretty severe performance hit, so it is vitally important that we utilize the fact that only 8 grid cells need to actually be checked. \ No newline at end of file diff --git a/boids.gif b/boids.gif new file mode 100644 index 0000000..60018f3 Binary files /dev/null and b/boids.gif differ diff --git a/graph1.png b/graph1.png new file mode 100644 index 0000000..6f25bbc Binary files /dev/null and b/graph1.png differ diff --git a/graph2.png b/graph2.png new file mode 100644 index 0000000..f208326 Binary files /dev/null and b/graph2.png differ diff --git a/graph3.png b/graph3.png new file mode 100644 index 0000000..be89550 Binary files /dev/null and b/graph3.png differ diff --git a/graph4.png b/graph4.png new file mode 100644 index 0000000..26bdfef Binary files /dev/null and b/graph4.png differ diff --git a/src/kernel.cu b/src/kernel.cu index aaf0fbf..4d7e424 100644 --- a/src/kernel.cu +++ b/src/kernel.cu @@ -5,6 +5,7 @@ #include #include "utilityCore.hpp" #include "kernel.h" +#include // LOOK-2.1 potentially useful for doing grid-based neighbor search #ifndef imax @@ -76,15 +77,15 @@ glm::vec3 *dev_vel2; // For efficient sorting and the uniform grid. These should always be parallel. int *dev_particleArrayIndices; // What index in dev_pos and dev_velX represents this particle? int *dev_particleGridIndices; // What grid cell is this particle in? -// needed for use with thrust -thrust::device_ptr dev_thrust_particleArrayIndices; -thrust::device_ptr dev_thrust_particleGridIndices; int *dev_gridCellStartIndices; // What part of dev_particleArrayIndices belongs int *dev_gridCellEndIndices; // to this cell? // TODO-2.3 - consider what additional buffers you might need to reshuffle // the position and velocity data to be coherent within cells. +// Parallel, memory coherent versions of the original simulation's dev_pos/vel1 +glm::vec3 *dev_pos_coherent; +glm::vec3 *dev_vel1_coherent; // LOOK-2.1 - Grid parameters based on simulation parameters. // These are automatically computed for you in Boids::initSimulation @@ -151,8 +152,14 @@ void Boids::initSimulation(int N) { cudaMalloc((void**)&dev_vel2, N * sizeof(glm::vec3)); checkCUDAErrorWithLine("cudaMalloc dev_vel2 failed!"); + cudaMalloc((void**)&dev_particleArrayIndices, N * sizeof(int)); + checkCUDAErrorWithLine("cudaMalloc dev_particleArrayIndices failed!"); + + cudaMalloc((void**)&dev_particleGridIndices, N * sizeof(int)); + checkCUDAErrorWithLine("cudaMalloc dev_particleGridIndices failed!"); + // LOOK-1.2 - This is a typical CUDA kernel invocation. - kernGenerateRandomPosArray<<>>(1, numObjects, + kernGenerateRandomPosArray << > > (1, numObjects, dev_pos, scene_scale); checkCUDAErrorWithLine("kernGenerateRandomPosArray failed!"); @@ -169,6 +176,19 @@ void Boids::initSimulation(int N) { gridMinimum.z -= halfGridWidth; // TODO-2.1 TODO-2.3 - Allocate additional buffers here. + + cudaMalloc((void**)&dev_pos_coherent, N * sizeof(glm::vec3)); + checkCUDAErrorWithLine("cudaMalloc dev_pos_copy failed!"); + + cudaMalloc((void**)&dev_vel1_coherent, N * sizeof(glm::vec3)); + checkCUDAErrorWithLine("cudaMalloc dev_vel1_copy failed!"); + + cudaMalloc((void**)&dev_gridCellStartIndices, gridCellCount * sizeof(int)); + checkCUDAErrorWithLine("cudaMalloc dev_gridCellStartIndices failed!"); + + cudaMalloc((void**)&dev_gridCellEndIndices, gridCellCount * sizeof(int)); + checkCUDAErrorWithLine("cudaMalloc dev_gridCellEndIndices failed!"); + cudaThreadSynchronize(); } @@ -210,8 +230,8 @@ __global__ void kernCopyVelocitiesToVBO(int N, glm::vec3 *vel, float *vbo, float void Boids::copyBoidsToVBO(float *vbodptr_positions, float *vbodptr_velocities) { dim3 fullBlocksPerGrid((numObjects + blockSize - 1) / blockSize); - kernCopyPositionsToVBO << > >(numObjects, dev_pos, vbodptr_positions, scene_scale); - kernCopyVelocitiesToVBO << > >(numObjects, dev_vel1, vbodptr_velocities, scene_scale); + kernCopyPositionsToVBO << > > (numObjects, dev_pos, vbodptr_positions, scene_scale); + kernCopyVelocitiesToVBO << > > (numObjects, dev_vel1, vbodptr_velocities, scene_scale); checkCUDAErrorWithLine("copyBoidsToVBO failed!"); @@ -230,10 +250,62 @@ void Boids::copyBoidsToVBO(float *vbodptr_positions, float *vbodptr_velocities) * in the `pos` and `vel` arrays. */ __device__ glm::vec3 computeVelocityChange(int N, int iSelf, const glm::vec3 *pos, const glm::vec3 *vel) { - // Rule 1: boids fly towards their local perceived center of mass, which excludes themselves - // Rule 2: boids try to stay a distance d away from each other - // Rule 3: boids try to match the speed of surrounding boids - return glm::vec3(0.0f, 0.0f, 0.0f); + glm::vec3 perceivedCenter(0); + glm::vec3 avoidanceVelocity(0); + glm::vec3 perceivedVelocity(0); + int r1_neighbors = 0; + int r2_neighbors = 0; + int r3_neighbors = 0; + glm::vec3 selfPos = pos[iSelf]; + + for (int i = 0; i < N; ++i) + { + glm::vec3 currentPos = pos[i]; + glm::vec3 currentDisplacement = currentPos - selfPos; + float displacementMagn = glm::length(currentDisplacement); + + // Rule 1: boids fly towards their local perceived center of mass, which excludes themselves + if (displacementMagn < rule1Distance) + { + perceivedCenter += currentPos; + r1_neighbors++; + } + + // Rule 2: boids try to stay a distance d away from each other + if (displacementMagn < rule2Distance) + { + avoidanceVelocity -= currentDisplacement; + r2_neighbors++; + } + + // Rule 3: boids try to match the speed of surrounding boids + if (displacementMagn < rule3Distance) + { + perceivedVelocity += vel[i]; + r3_neighbors++; + } + } + + if (r1_neighbors > 0) + { + perceivedCenter /= r1_neighbors; + } + + if (r3_neighbors > 0) + { + perceivedVelocity /= r3_neighbors; + } + + if (r1_neighbors + r2_neighbors + r3_neighbors > 0) + { + return (perceivedCenter - selfPos) * rule1Scale + + avoidanceVelocity * rule2Scale + + perceivedVelocity * rule3Scale; + } + else + { + return avoidanceVelocity * rule2Scale; + } } /** @@ -245,6 +317,20 @@ __global__ void kernUpdateVelocityBruteForce(int N, glm::vec3 *pos, // Compute a new velocity based on pos and vel1 // Clamp the speed // Record the new velocity into vel2. Question: why NOT vel1? + + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= N) { + return; + } + + glm::vec3 velChange = computeVelocityChange(N, index, pos, vel1); + + glm::vec3 newVel = vel1[index] + velChange; + if (glm::length(newVel) > maxSpeed) // clamp speed to max + { + newVel *= maxSpeed / glm::length(newVel); + } + vel2[index] = newVel; } /** @@ -278,17 +364,29 @@ __global__ void kernUpdatePos(int N, float dt, glm::vec3 *pos, glm::vec3 *vel) { // for(x) // for(y) // for(z)? Or some other order? -__device__ int gridIndex3Dto1D(int x, int y, int z, int gridResolution) { +/*__device__ int gridIndex3Dto1D(int x, int y, int z, int gridResolution) { return x + y * gridResolution + z * gridResolution * gridResolution; +}*/ + +__device__ int gridIndex3Dto1D(int x, int y, int z, int gridResolution) { + return z + y * gridResolution + x * gridResolution * gridResolution; } __global__ void kernComputeIndices(int N, int gridResolution, glm::vec3 gridMin, float inverseCellWidth, glm::vec3 *pos, int *indices, int *gridIndices) { - // TODO-2.1 - // - Label each boid with the index of its grid cell. - // - Set up a parallel array of integer indices as pointers to the actual - // boid data in pos and vel1/vel2 + // TODO-2.1 + // - Label each boid with the index of its grid cell. + // - Set up a parallel array of integer indices as pointers to the actual + // boid data in pos and vel1/vel2 + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= N) { + return; + } + glm::vec3 index3D = floor((pos[index] - gridMin) * inverseCellWidth); + int index1D = gridIndex3Dto1D(index3D.x, index3D.y, index3D.z, gridResolution); + gridIndices[index] = index1D; + indices[index] = index; } // LOOK-2.1 Consider how this could be useful for indicating that a cell @@ -300,12 +398,41 @@ __global__ void kernResetIntBuffer(int N, int *intBuffer, int value) { } } +__global__ void kernMakeArraysCoherent(int N, int *particleArrayIndices, + glm::vec3 *pos, glm::vec3 *pos_coherent, glm::vec3 *vel, glm::vec3 *vel_coherent) +{ + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= N) { + return; + } + pos_coherent[index] = pos[particleArrayIndices[index]]; + vel_coherent[index] = vel[particleArrayIndices[index]]; +} + __global__ void kernIdentifyCellStartEnd(int N, int *particleGridIndices, int *gridCellStartIndices, int *gridCellEndIndices) { // TODO-2.1 // Identify the start point of each cell in the gridIndices array. // This is basically a parallel unrolling of a loop that goes // "this index doesn't match the one before it, must be a new cell!" + int index = threadIdx.x + (blockIdx.x * blockDim.x); + int gridIdx = particleGridIndices[index]; + int gridIdxNext = particleGridIndices[index + 1]; + if (index == N - 1) + { + gridCellEndIndices[gridIdx] = index; + return; + } + + if (index > N - 1) { + return; + } + + if (gridIdx != gridIdxNext) + { + gridCellEndIndices[gridIdx] = index; + gridCellStartIndices[gridIdxNext] = index + 1; + } } __global__ void kernUpdateVelNeighborSearchScattered( @@ -322,6 +449,108 @@ __global__ void kernUpdateVelNeighborSearchScattered( // - Access each boid in the cell and compute velocity change from // the boids rules, if this boid is within the neighborhood distance. // - Clamp the speed change before putting the new speed in vel2 + + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= N) { + return; + } + + int actualBoidIndex = particleArrayIndices[index]; + glm::vec3 boidPos = pos[actualBoidIndex]; + glm::vec3 boidPosLocalToGrid = boidPos - gridMin; + glm::vec3 index3D = floor(boidPosLocalToGrid * inverseCellWidth); + glm::vec3 gridCellCenter = (index3D + glm::vec3(0.5f)) * cellWidth; + glm::vec3 gridOctant = glm::sign(boidPosLocalToGrid - gridCellCenter); + + glm::vec3 perceivedCenter(0); + glm::vec3 avoidanceVelocity(0); + glm::vec3 perceivedVelocity(0); + int r1_neighbors = 0; + int r2_neighbors = 0; + int r3_neighbors = 0; + + // Check all eight surrounding grid indices (with out of bounds checks) + for (int x = 0; x <= 1; ++x) + { + for (int y = 0; y <= 1; ++y) + { + for (int z = 0; z <= 1; ++z) + { + glm::vec3 currentGridIndex = glm::vec3(((float) x) * gridOctant.x, ((float) y) * gridOctant.y, ((float) z) * gridOctant.z); + currentGridIndex += index3D; + + if (((((int) currentGridIndex.x) >= 0 && ((int) currentGridIndex.x) < gridResolution) && + (((int) currentGridIndex.y) >= 0 && ((int) currentGridIndex.y) < gridResolution)) && + (((int) currentGridIndex.z) >= 0 && ((int) currentGridIndex.z) < gridResolution)) + { + int index1D = gridIndex3Dto1D(currentGridIndex.x, currentGridIndex.y, currentGridIndex.z, gridResolution); + for (int g = gridCellStartIndices[index1D]; g <= gridCellEndIndices[index1D]; ++g) + { + if (g < 0) break; + int currentBoid = particleArrayIndices[g]; + glm::vec3 currentBoidPos = pos[currentBoid]; + + // Boid + if (currentBoid != actualBoidIndex) + { + glm::vec3 currentDisplacement = currentBoidPos - boidPos; + float displacementMagn = glm::length(currentDisplacement); + + // Rule 1: boids fly towards their local perceived center of mass, which excludes themselves + if (displacementMagn < rule1Distance) + { + perceivedCenter += currentBoidPos; + r1_neighbors++; + } + + // Rule 2: boids try to stay a distance d away from each other + if (displacementMagn < rule2Distance) + { + avoidanceVelocity -= currentDisplacement; + r2_neighbors++; + } + + // Rule 3: boids try to match the speed of surrounding boids + if (displacementMagn < rule3Distance) + { + perceivedVelocity += vel1[currentBoid]; + r3_neighbors++; + } + } + } + } + } + } + } + + if (r1_neighbors > 0) + { + perceivedCenter /= r1_neighbors; + } + + if (r3_neighbors > 0) + { + perceivedVelocity /= r3_neighbors; + } + + glm::vec3 velChange(0); + if (r1_neighbors + r2_neighbors + r3_neighbors > 0) + { + velChange = (perceivedCenter - boidPos) * rule1Scale + + avoidanceVelocity * rule2Scale + + perceivedVelocity * rule3Scale; + } + else + { + velChange = avoidanceVelocity * rule2Scale; + } + + glm::vec3 newVel = vel1[actualBoidIndex] + velChange; + if (glm::length(newVel) > maxSpeed) // clamp speed to max + { + newVel *= maxSpeed / glm::length(newVel); + } + vel2[actualBoidIndex] = newVel; } __global__ void kernUpdateVelNeighborSearchCoherent( @@ -341,6 +570,114 @@ __global__ void kernUpdateVelNeighborSearchCoherent( // - Access each boid in the cell and compute velocity change from // the boids rules, if this boid is within the neighborhood distance. // - Clamp the speed change before putting the new speed in vel2 + + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= N) { + return; + } + + glm::vec3 boidPos = pos[index]; + glm::vec3 boidPosLocalToGrid = boidPos - gridMin; + glm::vec3 index3D = floor(boidPosLocalToGrid * inverseCellWidth); + glm::vec3 gridCellCenter = (index3D + glm::vec3(0.5f)) * cellWidth; + glm::vec3 gridOctant = glm::sign(boidPosLocalToGrid - gridCellCenter); + + glm::vec3 perceivedCenter(0); + glm::vec3 avoidanceVelocity(0); + glm::vec3 perceivedVelocity(0); + int r1_neighbors = 0; + int r2_neighbors = 0; + int r3_neighbors = 0; + + // Check all eight surrounding grid indices (with out of bounds checks) + for (int x = 0; x <= 1; ++x) + { + for (int y = 0; y <= 1; ++y) + { + for (int z = 0; z <= 1; ++z) + { + glm::vec3 currentGridIndex = glm::vec3(((float)x) * gridOctant.x, ((float)y) * gridOctant.y, ((float)z) * gridOctant.z); + currentGridIndex += index3D; + + if (((((int)currentGridIndex.x) >= 0 && ((int)currentGridIndex.x) < gridResolution) && + (((int)currentGridIndex.y) >= 0 && ((int)currentGridIndex.y) < gridResolution)) && + (((int)currentGridIndex.z) >= 0 && ((int)currentGridIndex.z) < gridResolution)) + { + int index1D = gridIndex3Dto1D(currentGridIndex.x, currentGridIndex.y, currentGridIndex.z, gridResolution); + for (int g = gridCellStartIndices[index1D]; g <= gridCellEndIndices[index1D]; ++g) + { + if (g < 0) break; + glm::vec3 currentBoidPos = pos[g]; + + // Boid + if (g != index) + { + glm::vec3 currentDisplacement = currentBoidPos - boidPos; + float displacementMagn = glm::length(currentDisplacement); + + // Rule 1: boids fly towards their local perceived center of mass, which excludes themselves + if (displacementMagn < rule1Distance) + { + perceivedCenter += currentBoidPos; + r1_neighbors++; + } + + // Rule 2: boids try to stay a distance d away from each other + if (displacementMagn < rule2Distance) + { + avoidanceVelocity -= currentDisplacement; + r2_neighbors++; + } + + // Rule 3: boids try to match the speed of surrounding boids + if (displacementMagn < rule3Distance) + { + perceivedVelocity += vel1[g]; + r3_neighbors++; + } + } + } + } + } + } + } + + if (r1_neighbors > 0) + { + perceivedCenter /= r1_neighbors; + } + + if (r3_neighbors > 0) + { + perceivedVelocity /= r3_neighbors; + } + + glm::vec3 velChange(0); + if (r1_neighbors + r2_neighbors + r3_neighbors > 0) + { + velChange = (perceivedCenter - boidPos) * rule1Scale + + avoidanceVelocity * rule2Scale + + perceivedVelocity * rule3Scale; + } + else + { + velChange = avoidanceVelocity * rule2Scale; + } + + glm::vec3 newVel = vel1[index] + velChange; + if (glm::length(newVel) > maxSpeed) // clamp speed to max + { + newVel *= maxSpeed / glm::length(newVel); + } + vel2[index] = newVel; +} + +__global__ void kernPingPongBuffers(int N, glm::vec3 *b1, glm::vec3 *b2) { + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= N) { + return; + } + b1[index] = b2[index]; } /** @@ -348,7 +685,11 @@ __global__ void kernUpdateVelNeighborSearchCoherent( */ void Boids::stepSimulationNaive(float dt) { // TODO-1.2 - use the kernels you wrote to step the simulation forward in time. - // TODO-1.2 ping-pong the velocity buffers + // TODO-1.2 ping-pong the velocity buffers. + dim3 fullBlocksPerGrid((numObjects + blockSize - 1) / blockSize); + kernUpdateVelocityBruteForce << > > (numObjects, dev_pos, dev_vel1, dev_vel2); + kernUpdatePos << > > (numObjects, dt, dev_pos, dev_vel2); + kernPingPongBuffers << > > (numObjects, dev_vel1, dev_vel2); } void Boids::stepSimulationScatteredGrid(float dt) { @@ -364,6 +705,26 @@ void Boids::stepSimulationScatteredGrid(float dt) { // - Perform velocity updates using neighbor search // - Update positions // - Ping-pong buffers as needed + + dim3 fullBlocksPerGrid((numObjects + blockSize - 1) / blockSize); + dim3 fullBlocksPerGrid_GridCell((gridCellCount + blockSize - 1) / blockSize); + + // Initialize the grid index start/cell array to -1 + kernResetIntBuffer << > > (gridCellCount, dev_gridCellStartIndices, -1); + kernResetIntBuffer << > > (gridCellCount, dev_gridCellEndIndices, -1); + + kernComputeIndices << > > (numObjects, gridSideCount, gridMinimum, gridInverseCellWidth, dev_pos, dev_particleArrayIndices, dev_particleGridIndices); + + // Sort with thrust + thrust::device_ptr dev_thrust_grid_indices(dev_particleGridIndices); + thrust::device_ptr dev_thrust_array_indices(dev_particleArrayIndices); + thrust::sort_by_key(dev_thrust_grid_indices, dev_thrust_grid_indices + numObjects, dev_thrust_array_indices); + + kernIdentifyCellStartEnd << > > (numObjects, dev_particleGridIndices, dev_gridCellStartIndices, dev_gridCellEndIndices); + kernUpdateVelNeighborSearchScattered << > > (numObjects, gridSideCount, gridMinimum, gridInverseCellWidth, gridCellWidth, dev_gridCellStartIndices, dev_gridCellEndIndices, dev_particleArrayIndices, dev_pos, dev_vel1, dev_vel2); + + kernUpdatePos << > > (numObjects, dt, dev_pos, dev_vel2); + kernPingPongBuffers << > > (numObjects, dev_vel1, dev_vel2); } void Boids::stepSimulationCoherentGrid(float dt) { @@ -382,12 +743,40 @@ void Boids::stepSimulationCoherentGrid(float dt) { // - Perform velocity updates using neighbor search // - Update positions // - Ping-pong buffers as needed. THIS MAY BE DIFFERENT FROM BEFORE. + + dim3 fullBlocksPerGrid((numObjects + blockSize - 1) / blockSize); + dim3 fullBlocksPerGrid_GridCell((gridCellCount + blockSize - 1) / blockSize); + + // Initialize the grid index start/cell array to -1 + kernResetIntBuffer << > > (gridCellCount, dev_gridCellStartIndices, -1); + kernResetIntBuffer << > > (gridCellCount, dev_gridCellEndIndices, -1); + + kernComputeIndices << > > (numObjects, gridSideCount, gridMinimum, gridInverseCellWidth, dev_pos, dev_particleArrayIndices, dev_particleGridIndices); + + // Sort with thrust + thrust::device_ptr dev_thrust_grid_indices(dev_particleGridIndices); + thrust::device_ptr dev_thrust_array_indices(dev_particleArrayIndices); + thrust::sort_by_key(dev_thrust_grid_indices, dev_thrust_grid_indices + numObjects, dev_thrust_array_indices); + + kernMakeArraysCoherent << > > (numObjects, dev_particleArrayIndices, dev_pos, dev_pos_coherent, dev_vel1, dev_vel1_coherent); + kernIdentifyCellStartEnd << > > (numObjects, dev_particleGridIndices, dev_gridCellStartIndices, dev_gridCellEndIndices); + kernUpdateVelNeighborSearchCoherent << > > (numObjects, gridSideCount, gridMinimum, gridInverseCellWidth, gridCellWidth, dev_gridCellStartIndices, dev_gridCellEndIndices, dev_pos_coherent, dev_vel1_coherent, dev_vel2); + + kernPingPongBuffers << > > (numObjects, dev_pos, dev_pos_coherent); + kernUpdatePos << > > (numObjects, dt, dev_pos, dev_vel2); + kernPingPongBuffers << > > (numObjects, dev_vel1, dev_vel2); } void Boids::endSimulation() { cudaFree(dev_vel1); cudaFree(dev_vel2); cudaFree(dev_pos); + cudaFree(dev_particleArrayIndices); + cudaFree(dev_particleGridIndices); + cudaFree(dev_gridCellStartIndices); + cudaFree(dev_gridCellEndIndices); + cudaFree(dev_pos_coherent); + cudaFree(dev_vel1_coherent); // TODO-2.1 TODO-2.3 - Free any additional buffers here. } diff --git a/src/main.cpp b/src/main.cpp index a29471d..b2a18fc 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -14,11 +14,11 @@ // LOOK-2.1 LOOK-2.3 - toggles for UNIFORM_GRID and COHERENT_GRID #define VISUALIZE 1 -#define UNIFORM_GRID 0 -#define COHERENT_GRID 0 +#define UNIFORM_GRID 1 +#define COHERENT_GRID 1 // LOOK-1.2 - change this to adjust particle count in the simulation -const int N_FOR_VIS = 5000; +const int N_FOR_VIS = 10000; const float DT = 0.2f; /**