diff --git a/README.md b/README.md index ee39093..f68fae1 100644 --- a/README.md +++ b/README.md @@ -1,11 +1,51 @@ **University of Pennsylvania, CIS 5650: GPU Programming and Architecture, Project 1 - Flocking** -* (TODO) YOUR NAME HERE - * (TODO) [LinkedIn](), [personal website](), [twitter](), etc. -* Tested on: (TODO) Windows 22, i7-2222 @ 2.22GHz 22GB, GTX 222 222MB (Moore 2222 Lab) +* Joanne Li + * [LinkedIn](https://www.linkedin.com/in/zhuoran-li-856658244/) +* Tested on: Windows 11, AMD Ryzen 5 5600H @ 3.30 GHz 16.0GB, NVIDIA GeForce RTX 3050 Laptop GPU 4GB -### (TODO: Your 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.) +## Final results +50000 boids: +![](images/50k_boids_gif.gif) +![](images/50k_boids.png) + + +5000 boids: +![](images/5k_boids_gif.gif) +![](images/5k_boids.png) + + +## Performance analysis +Frame rate with visualization: +![](images/graph1.png) +Frame rate without visualization: +![](images/graph2.png) + +* For each implementation, how does changing the number of boids affect performance? Why do you think this is? + + **Naive implementation**: When I increase the number of boids, performance drops a lot. This makes sense because every boid checks against all other boids, so the work grows like N^2. + + **Uniform grid (scattered and coherent)**: Performance also gets worse as I add more boids, but not as dramatically. Each boid only checks neighbors in a few cells instead of all boids, so the growth is closer to linear. The grid helps cut down the comparisons. + + +* For each implementation, how does changing the block count and block size affect performance? Why do you think this is? + + If the block size is too small, the GPU isn't fully used, so performance suffers. If the block size is too big (more than what the hardware can handle well), some resources are wasted. + + Changing the number of blocks mostly changes how work is split; as long as there are enough blocks to keep all GPU cores busy, performance is stable. + + +* For the coherent uniform grid: did you experience any performance improvements with the more coherent uniform grid? Was this the outcome you expected? Why or why not? + + Yes, the coherent grid was faster than the scattered one. This is because, after sorting, positions and velocities of boids in the same cell are stored next to each other in memory. That means memory accesses are more friendly, so the GPU can grab data efficiently. + + I expected this improvement, since the main difference between scattered vs coherent is how memory is accessed. GPUs are very sensitive to memory patterns, so coherent layout helps. + + +* Did changing cell width and checking 27 vs 8 neighboring cells affect performance? Why or why not? + + Yes, changing the cell width affects performance. If the cell width is too small, there are many empty cells, and each boid still needs to check lots of cells, which wastes time. On the other hand, If the cell width is too big, too many boids fall into the same cell, so each boid ends up checking way more neighbors than needed. + + For 27 vs 8 cells: it's not always true that 27 is slower just because it's more cells. Sometimes checking 27 cells is better if the cell width is small, because each cell has very few boids, so the total number of comparisons stays low. On the other hand, with larger cells, 8 is usually enough, because most neighbors are already in those few cells. \ No newline at end of file diff --git a/images/50k_boids.png b/images/50k_boids.png new file mode 100644 index 0000000..2384087 Binary files /dev/null and b/images/50k_boids.png differ diff --git a/images/50k_boids_gif.gif b/images/50k_boids_gif.gif new file mode 100644 index 0000000..8913498 Binary files /dev/null and b/images/50k_boids_gif.gif differ diff --git a/images/5k_boids.png b/images/5k_boids.png new file mode 100644 index 0000000..fa02ee1 Binary files /dev/null and b/images/5k_boids.png differ diff --git a/images/5k_boids_gif.gif b/images/5k_boids_gif.gif new file mode 100644 index 0000000..7f58350 Binary files /dev/null and b/images/5k_boids_gif.gif differ diff --git a/images/graph1.png b/images/graph1.png new file mode 100644 index 0000000..c710573 Binary files /dev/null and b/images/graph1.png differ diff --git a/images/graph2.png b/images/graph2.png new file mode 100644 index 0000000..c3a69e5 Binary files /dev/null and b/images/graph2.png differ diff --git a/reports/performance_analysis.xlsx b/reports/performance_analysis.xlsx new file mode 100644 index 0000000..c9fcede Binary files /dev/null and b/reports/performance_analysis.xlsx differ diff --git a/src/kernel.cu b/src/kernel.cu index 7149917..e48e28d 100644 --- a/src/kernel.cu +++ b/src/kernel.cu @@ -52,7 +52,7 @@ void checkCUDAError(const char *msg, int line = -1) { // LOOK-1.2 Parameters for the boids algorithm. // These worked well in our reference implementation. #define rule1Distance 5.0f -#define rule2Distance 3.0f +#define rule2Distance 2.0f #define rule3Distance 5.0f #define rule1Scale 0.01f @@ -96,6 +96,9 @@ 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. +glm::vec3* dev_newPos; +glm::vec3* dev_newVel; + // LOOK-2.1 - Grid parameters based on simulation parameters. // These are automatically computed for you in Boids::initSimulation int gridCellCount; @@ -179,6 +182,27 @@ void Boids::initSimulation(int N) { gridMinimum.z -= halfGridWidth; // TODO-2.1 TODO-2.3 - Allocate additional buffers here. + cudaMalloc((void**)&dev_particleArrayIndices, numObjects * sizeof(int)); + checkCUDAErrorWithLine("cudaMalloc dev_particleArrayIndices failed!"); + + cudaMalloc((void**)&dev_particleGridIndices, numObjects * sizeof(int)); + checkCUDAErrorWithLine("cudaMalloc dev_particleGridIndices 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!"); + + dev_thrust_particleArrayIndices = thrust::device_ptr(dev_particleArrayIndices); + dev_thrust_particleGridIndices = thrust::device_ptr(dev_particleGridIndices); + + cudaMalloc((void**)&dev_newPos, N * sizeof(glm::vec3)); + checkCUDAErrorWithLine("cudaMalloc dev_newPos failed!"); + + cudaMalloc((void**)&dev_newVel, N * sizeof(glm::vec3)); + checkCUDAErrorWithLine("cudaMalloc dev_newVel failed!"); + cudaDeviceSynchronize(); } @@ -239,11 +263,53 @@ void Boids::copyBoidsToVBO(float *vbodptr_positions, float *vbodptr_velocities) * Compute the new velocity on the body with index `iSelf` due to the `N` boids * 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); +__device__ glm::vec3 computeVelocityChange(int N, int iSelf, const glm::vec3* pos, const glm::vec3* vel) { + glm::vec3 velocity(0.f); + int rule1Nb = 0; + int rule3Nb = 0; + glm::vec3 pc(0.f); + glm::vec3 c(0.f); + glm::vec3 pv(0.f); + + for (int i = 0; i < N; i++) { + if (i == iSelf) continue; // Check if the boid is itself + glm::vec3 diff = pos[i] - pos[iSelf]; + + // Rule 1: boids fly towards their local perceived center of mass, which excludes themselves + if (dot(diff, diff) < rule1Distance * rule1Distance) { // Check if the boid is inside the given distance + pc += pos[i]; + rule1Nb++; + } + + // Rule 2: boids try to stay a distance d away from each other + if (dot(diff, diff) < rule2Distance * rule2Distance) { + c -= pos[i] - pos[iSelf]; + } + + // Rule 3: boids try to match the speed of surrounding boids + if (dot(diff, diff) < rule3Distance * rule3Distance) { + pv += vel[i]; + rule3Nb++; + } + } + + // Rule 1 + pc = (rule1Nb == 0) ? pos[iSelf] : (pc / (float)rule1Nb); // Handle the case that the boid itself has no neibours + velocity += (pc - pos[iSelf]) * rule1Scale; + + // Rule 2 + velocity += c * rule2Scale; + + // Rule 3 + pv = (rule3Nb == 0) ? glm::vec3(0.f) : (pv / (float)rule3Nb); + velocity += pv * rule3Scale; + + velocity += vel[iSelf]; + if (glm::length(velocity) > maxSpeed) { + velocity = glm::normalize(velocity) * maxSpeed; + } + + return velocity; } /** @@ -251,10 +317,12 @@ __device__ glm::vec3 computeVelocityChange(int N, int iSelf, const glm::vec3 *po * For each of the `N` bodies, update its position based on its current velocity. */ __global__ void kernUpdateVelocityBruteForce(int N, glm::vec3 *pos, - glm::vec3 *vel1, glm::vec3 *vel2) { - // Compute a new velocity based on pos and vel1 - // Clamp the speed - // Record the new velocity into vel2. Question: why NOT vel1? + glm::vec3 *vel1, glm::vec3 *vel2) { + // Compute a new velocity based on pos and vel1 + // Clamp the speed + // Record the new velocity into vel2. Question: why NOT vel1? + int idx = blockIdx.x * blockDim.x + threadIdx.x; + vel2[idx] = computeVelocityChange(N, idx, pos, vel1); } /** @@ -299,6 +367,11 @@ __global__ void kernComputeIndices(int N, int gridResolution, // - 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 = (blockIdx.x * blockDim.x) + threadIdx.x; + if (index >= N) return; + glm::vec3 gridIndex = floor(pos[index] - gridMin) * inverseCellWidth; + gridIndices[index] = gridIndex3Dto1D(gridIndex.x, gridIndex.y, gridIndex.z, gridResolution); + indices[index] = index; } // LOOK-2.1 Consider how this could be useful for indicating that a cell @@ -316,6 +389,29 @@ __global__ void kernIdentifyCellStartEnd(int N, int *particleGridIndices, // 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 = (blockIdx.x * blockDim.x) + threadIdx.x; + if (index >= N) return; + + int gridIdx = particleGridIndices[index]; + + // First + if (index == 0) { + gridCellStartIndices[gridIdx] = 0; + } + + // Middle + if (index != 0) { + int gridIdxPrev = particleGridIndices[index - 1]; + if (gridIdx != gridIdxPrev) { + gridCellStartIndices[gridIdx] = index; + gridCellEndIndices[gridIdxPrev] = index - 1; + } + } + + // Last + if (index == N - 1) { + gridCellEndIndices[gridIdx] = N - 1; + } } __global__ void kernUpdateVelNeighborSearchScattered( @@ -327,11 +423,90 @@ __global__ void kernUpdateVelNeighborSearchScattered( // TODO-2.1 - Update a boid's velocity using the uniform grid to reduce // the number of boids that need to be checked. // - Identify the grid cell that this particle is in + + int idx = (blockIdx.x * blockDim.x) + threadIdx.x; + if (idx >= N) return; + glm::vec3 pPos = pos[idx]; + glm::vec3 gridIdx3D = (pPos - gridMin) * inverseCellWidth; + int gridIdx = gridIndex3Dto1D(gridIdx3D.x, gridIdx3D.y, gridIdx3D.z, gridResolution); + // - Identify which cells may contain neighbors. This isn't always 8. - // - For each cell, read the start/end indices in the boid pointer array. - // - Access each boid in the cell and compute velocity change from - // the boids rules, if this boid is within the neighborhood distance. + float r = imax(imax(rule1Distance, rule2Distance), rule3Distance); + + int minX = floor((pPos.x - r - gridMin.x) * inverseCellWidth); + int maxX = floor((pPos.x + r - gridMin.x) * inverseCellWidth); + + int minY = floor((pPos.y - r - gridMin.y) * inverseCellWidth); + int maxY = floor((pPos.y + r - gridMin.y) * inverseCellWidth); + + int minZ = floor((pPos.z - r - gridMin.z) * inverseCellWidth); + int maxZ = floor((pPos.z + r - gridMin.z) * inverseCellWidth); + + + glm::vec3 velocity(0.f); + int rule1Nb = 0; + int rule3Nb = 0; + glm::vec3 pc(0.f); + glm::vec3 c(0.f); + glm::vec3 pv(0.f); + + // - For each cell, read the start/end indices in the boid pointer array. + for (int i = minX; i < maxX + 1; ++i) { + for (int j = minY; j < maxY + 1; ++j) { + for (int k = minZ; k < maxZ + 1; ++k) { + int grid = gridIndex3Dto1D(i, j, k, gridResolution); + + int start = gridCellStartIndices[grid]; + if (start == -1) continue; // The grid contains no particles + int end = gridCellEndIndices[grid]; + + // - Access each boid in the cell and compute velocity change from + // the boids rules, if this boid is within the neighborhood distance. + for (int p = start; p <= end; ++p) { + int pIdx = particleArrayIndices[p]; + if (pIdx == idx) continue; + + glm::vec3 diff = pos[pIdx] - pos[idx]; + + // Rule 1 + if (dot(diff, diff) < rule1Distance * rule1Distance) { + pc += pos[pIdx]; + rule1Nb++; + } + + // Rule 2 + if (dot(diff, diff) < rule2Distance * rule2Distance) { + c -= pos[pIdx] - pos[idx]; + } + + // Rule 3 + if (dot(diff, diff) < rule3Distance * rule3Distance) { + pv += vel1[pIdx]; + rule3Nb++; + } + } + } + } + } + + // Rule 1 + pc = (rule1Nb == 0) ? pos[idx] : (pc / (float)rule1Nb); + velocity += (pc - pos[idx]) * rule1Scale; + + // Rule 2 + velocity += c * rule2Scale; + + // Rule 3 + pv = (rule3Nb == 0) ? glm::vec3(0.f) : (pv / (float)rule3Nb); + velocity += pv * rule3Scale; + + velocity += vel1[idx]; + if (glm::length(velocity) > maxSpeed) { + velocity = glm::normalize(velocity) * maxSpeed; + } + // - Clamp the speed change before putting the new speed in vel2 + vel2[idx] = velocity; } __global__ void kernUpdateVelNeighborSearchCoherent( @@ -343,55 +518,209 @@ __global__ void kernUpdateVelNeighborSearchCoherent( // except with one less level of indirection. // This should expect gridCellStartIndices and gridCellEndIndices to refer // directly to pos and vel1. + // - Identify the grid cell that this particle is in - // - Identify which cells may contain neighbors. This isn't always 8. - // - For each cell, read the start/end indices in the boid pointer array. - // DIFFERENCE: For best results, consider what order the cells should be - // checked in to maximize the memory benefits of reordering the boids data. - // - 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 idx = (blockIdx.x * blockDim.x) + threadIdx.x; + if (idx >= N) return; + glm::vec3 pPos = pos[idx]; + glm::vec3 gridIdx3D = (pPos - gridMin) * inverseCellWidth; + int gridIdx = gridIndex3Dto1D(gridIdx3D.x, gridIdx3D.y, gridIdx3D.z, gridResolution); + + // - Identify which cells may contain neighbors. This isn't always 8. + float r = imax(imax(rule1Distance, rule2Distance), rule3Distance); + + int minX = floor((pPos.x - r - gridMin.x) * inverseCellWidth); + int maxX = floor((pPos.x + r - gridMin.x) * inverseCellWidth); + + int minY = floor((pPos.y - r - gridMin.y) * inverseCellWidth); + int maxY = floor((pPos.y + r - gridMin.y) * inverseCellWidth); + + int minZ = floor((pPos.z - r - gridMin.z) * inverseCellWidth); + int maxZ = floor((pPos.z + r - gridMin.z) * inverseCellWidth); + + glm::vec3 velocity(0.f); + int rule1Nb = 0; + int rule3Nb = 0; + glm::vec3 pc(0.f); + glm::vec3 c(0.f); + glm::vec3 pv(0.f); + + // - For each cell, read the start/end indices in the boid pointer array. + // DIFFERENCE: For best results, consider what order the cells should be + // checked in to maximize the memory benefits of reordering the boids data. + for (int k = minZ; k < maxZ + 1; ++k) { + for (int j = minY; j < maxY + 1; ++j) { + for (int i = minX; i < maxX + 1; ++i) { + int grid = gridIndex3Dto1D(i, j, k, gridResolution); + + int start = gridCellStartIndices[grid]; + if (start == -1) continue; // The grid contains no particles + int end = gridCellEndIndices[grid]; + + // - Access each boid in the cell and compute velocity change from + // the boids rules, if this boid is within the neighborhood distance. + for (int p = start; p <= end; ++p) { + if (p == idx) continue; + + glm::vec3 diff = pos[p] - pos[idx]; + + // Rule 1 + if (dot(diff, diff) < rule1Distance * rule1Distance) { + pc += pos[p]; + rule1Nb++; + } + + // Rule 2 + if (dot(diff, diff) < rule2Distance * rule2Distance) { + c -= pos[p] - pos[idx]; + } + + // Rule 3 + if (dot(diff, diff) < rule3Distance * rule3Distance) { + pv += vel1[p]; + rule3Nb++; + } + } + } + } + } + + // Rule 1 + pc = (rule1Nb == 0) ? pos[idx] : (pc / (float)rule1Nb); + velocity += (pc - pos[idx]) * rule1Scale; + + // Rule 2 + velocity += c * rule2Scale; + + // Rule 3 + pv = (rule3Nb == 0) ? glm::vec3(0.f) : (pv / (float)rule3Nb); + velocity += pv * rule3Scale; + + velocity += vel1[idx]; + if (glm::length(velocity) > maxSpeed) { + velocity = glm::normalize(velocity) * maxSpeed; + } + + // - Clamp the speed change before putting the new speed in vel2 + vel2[idx] = velocity; } +__global__ void kernShufflePosVel(int N, int* particleArrayIndices, + glm::vec3* pos, glm::vec3* vel, + glm::vec3* newPos, glm::vec3* newVel) { + + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= N) return; + + int origIdx = particleArrayIndices[index]; + + newPos[index] = pos[origIdx]; + newVel[index] = vel[origIdx]; +} + + /** * Step the entire N-body simulation by `dt` seconds. */ 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 + int blocks = (numObjects + blockSize - 1) / blockSize; + kernUpdateVelocityBruteForce<<>>(numObjects, dev_pos, dev_vel1, dev_vel2); + kernUpdatePos<<>>(numObjects, dt, dev_pos, dev_vel2); + dev_vel1 = dev_vel2; } void Boids::stepSimulationScatteredGrid(float dt) { // TODO-2.1 // Uniform Grid Neighbor search using Thrust sort. + + int blocks = (numObjects + blockSize - 1) / blockSize; + int blocksForCell = (gridCellCount + blockSize - 1) / blockSize; + // In Parallel: // - label each particle with its array index as well as its grid index. // Use 2x width grids. + kernComputeIndices<<>>(numObjects, gridSideCount, gridMinimum, gridInverseCellWidth, + dev_pos, dev_particleArrayIndices, dev_particleGridIndices); + checkCUDAErrorWithLine("kernComputeIndices failed!"); + // - Unstable key sort using Thrust. A stable sort isn't necessary, but you // are welcome to do a performance comparison. + thrust::sort_by_key(dev_thrust_particleGridIndices, dev_thrust_particleGridIndices + numObjects, dev_thrust_particleArrayIndices); + // - Naively unroll the loop for finding the start and end indices of each // cell's data pointers in the array of boid indices + kernResetIntBuffer<<>>(gridCellCount, dev_gridCellStartIndices, -1); + kernResetIntBuffer<<>>(gridCellCount, dev_gridCellEndIndices, -1); + checkCUDAErrorWithLine("kernResetIntBuffer failed!"); + + kernIdentifyCellStartEnd<<>>(numObjects, dev_particleGridIndices, + dev_gridCellStartIndices, dev_gridCellEndIndices); + checkCUDAErrorWithLine("kernIdentityCellStartEnd failed!"); + // - Perform velocity updates using neighbor search + kernUpdateVelNeighborSearchScattered<<>>(numObjects, gridSideCount, gridMinimum, gridInverseCellWidth, + gridCellWidth, dev_gridCellStartIndices, dev_gridCellEndIndices, + dev_particleArrayIndices, dev_pos, dev_vel1, dev_vel2); + checkCUDAErrorWithLine("kernUpdateVelNeighborSearchScattered failed!"); + // - Update positions + kernUpdatePos<<>>(numObjects, dt, dev_pos, dev_vel2); + checkCUDAErrorWithLine("kernUpdatePos failed!"); + // - Ping-pong buffers as needed + dev_vel1 = dev_vel2; } void Boids::stepSimulationCoherentGrid(float dt) { // TODO-2.3 - start by copying Boids::stepSimulationNaiveGrid // Uniform Grid Neighbor search using Thrust sort on cell-coherent data. + + int blocks = (numObjects + blockSize - 1) / blockSize; + int blocksForCell = (gridCellCount + blockSize - 1) / blockSize; + // In Parallel: - // - Label each particle with its array index as well as its grid index. - // Use 2x width grids + // - label each particle with its array index as well as its grid index. + // Use 2x width grids. + kernComputeIndices<<>>(numObjects, gridSideCount, gridMinimum, gridInverseCellWidth, + dev_pos, dev_particleArrayIndices, dev_particleGridIndices); + checkCUDAErrorWithLine("kernComputeIndices failed!"); + // - Unstable key sort using Thrust. A stable sort isn't necessary, but you // are welcome to do a performance comparison. - // - Naively unroll the loop for finding the start and end indices of each - // cell's data pointers in the array of boid indices + thrust::sort_by_key(dev_thrust_particleGridIndices, dev_thrust_particleGridIndices + numObjects, dev_thrust_particleArrayIndices); + // - BIG DIFFERENCE: use the rearranged array index buffer to reshuffle all // the particle data in the simulation array. // CONSIDER WHAT ADDITIONAL BUFFERS YOU NEED + kernShufflePosVel<<>>(numObjects, dev_particleArrayIndices, dev_pos, dev_vel1, + dev_newPos, dev_newVel); + checkCUDAErrorWithLine("kernShufflePosVel failed!"); + + // - Naively unroll the loop for finding the start and end indices of each + // cell's data pointers in the array of boid indices + kernResetIntBuffer<<>>(gridCellCount, dev_gridCellStartIndices, -1); + kernResetIntBuffer<<>>(gridCellCount, dev_gridCellEndIndices, -1); + checkCUDAErrorWithLine("kernResetIntBuffer failed!"); + + kernIdentifyCellStartEnd<<>>(numObjects, dev_particleGridIndices, + dev_gridCellStartIndices, dev_gridCellEndIndices); + checkCUDAErrorWithLine("kernIdentityCellStartEnd failed!"); + // - Perform velocity updates using neighbor search - // - Update positions - // - Ping-pong buffers as needed. THIS MAY BE DIFFERENT FROM BEFORE. + kernUpdateVelNeighborSearchCoherent<<>>(numObjects, gridSideCount, gridMinimum, + gridInverseCellWidth, gridCellWidth, dev_gridCellStartIndices, dev_gridCellEndIndices, + dev_newPos, dev_newVel, dev_vel2); + checkCUDAErrorWithLine("kernUpdateVelNeighborSearchCoherent failed!"); + + // Update positions + kernUpdatePos<<>>(numObjects, dt, dev_newPos, dev_vel2); + checkCUDAErrorWithLine("kernUpdatePos failed!"); + + // Ping-pong buffers + dev_pos = dev_newPos; + dev_vel1 = dev_vel2; } void Boids::endSimulation() { @@ -400,6 +729,13 @@ void Boids::endSimulation() { cudaFree(dev_pos); // TODO-2.1 TODO-2.3 - Free any additional buffers here. + cudaFree(dev_particleArrayIndices); + cudaFree(dev_particleGridIndices); + cudaFree(dev_gridCellStartIndices); + cudaFree(dev_gridCellEndIndices); + + cudaFree(dev_newPos); + cudaFree(dev_newVel); } void Boids::unitTest() { diff --git a/src/main.cpp b/src/main.cpp index 9c917c0..2fa4a53 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -23,7 +23,7 @@ // LOOK-2.1 LOOK-2.3 - toggles for UNIFORM_GRID and COHERENT_GRID #define VISUALIZE 1 -#define UNIFORM_GRID 0 +#define UNIFORM_GRID 1 #define COHERENT_GRID 0 // LOOK-1.2 - change this to adjust particle count in the simulation