diff --git a/README.md b/README.md index ee39093..1770fe6 100644 --- a/README.md +++ b/README.md @@ -1,11 +1,33 @@ **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) +* Kyle Bauer + * [LinkedIn](https://www.linkedin.com/in/kyle-bauer-75bb25171/), [twitter](https://x.com/KyleBauer414346) +* Tested on: Windows 10, i-7 12700 @ 2.1GHz 32GB, NVIDIA T1000 4GB (CETS Virtual Lab) -### (TODO: Your README) +![](images/boids.gif) -Include screenshots, analysis, etc. (Remember, this is public, so don't put -anything here that you don't want to share with the world.) +![](images/screencap_1.png) + +![](images/screencap_2.png) + +8 Neighbors | 27 Neighbors +:-------------------------:|:-------------------------: +![](images/128x8.png) | ![](images/128x27.png) +![](images/256x8.png) | ![](images/256x27.png) + +### Q1 For each implementation, how does changing the number of boids affect performance? Why do you think this is? + +While the performance noticeably decreased each time the boid count was doubled, the way each implementation changed varied greatly. The naive implementation seemed to drop a third of its framerate each double, whereas the scattered uniform grid usually did not even lose half of its framerate. The coherent uniform grid performed the best out of all three falling off even more gradually than the scattered variant. The naive implementation falling off the fastest makes the most sense as adding a new boid requires one additional check for every other boid in addition to the new boid having to check all the existing boids, quickly creating more work for each boid as the number of boids increases. This is different compared to the scattered and coherent grids as those implementations have culled most of the boids around them. However, those two implementations still see noticeable performance drop-offs because they are still affected by densly packed boids - Every close boids still need to accounted for when figuring out how each boid reacts. + +### Q2 For each implementation, how does changing the block count and block size affect performance? Why do you think this is? + +Changing the block size and block count did not noticeably affect performance. Due to the implementation, block size and block count are linked in an attempt to fill the grid as best as possible. Therefore, changing the block size will modify the block count and they should approximately balance eachother out. + +### Q3 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? + +The coherent grid performed noticeably better than the scattered grid which I found to be unexpected. I wasn't sure if the extra work required to make the grid coherent would be worth the performance gain from having better locality in memory, but clearly by the performance analysis it outperformed the scattered grid. + +### Q4 Did changing cell width and checking 27 vs 8 neighboring cells affect performance? Why or why not? Be careful: it is insufficient (and possibly incorrect) to say that 27-cell is slower simply because there are more cells to check! + +There was a significant increase in performance when checking 27 neighboring cells versus checking just 8 cells. This was most noticable for the dense (40k boids) coherent grid. I believe this is happening because the more densely packed grids benefit from a smaller cell width - It allows them to cull more boids early on. At those denser scales, the culling of boids out-weighed the performance hit of organizing a finer coherent grid. diff --git a/images/128x27.png b/images/128x27.png new file mode 100644 index 0000000..2de4bf3 Binary files /dev/null and b/images/128x27.png differ diff --git a/images/128x8.png b/images/128x8.png new file mode 100644 index 0000000..656c333 Binary files /dev/null and b/images/128x8.png differ diff --git a/images/256x27.png b/images/256x27.png new file mode 100644 index 0000000..35838b6 Binary files /dev/null and b/images/256x27.png differ diff --git a/images/256x8.png b/images/256x8.png new file mode 100644 index 0000000..502ca61 Binary files /dev/null and b/images/256x8.png differ diff --git a/images/boids.gif b/images/boids.gif new file mode 100644 index 0000000..f294075 Binary files /dev/null and b/images/boids.gif differ diff --git a/images/screencap_1.png b/images/screencap_1.png new file mode 100644 index 0000000..6b0116d Binary files /dev/null and b/images/screencap_1.png differ diff --git a/images/screencap_2.png b/images/screencap_2.png new file mode 100644 index 0000000..5c88baa Binary files /dev/null and b/images/screencap_2.png differ diff --git a/src/kernel.cu b/src/kernel.cu index 74dffcb..ff86462 100644 --- a/src/kernel.cu +++ b/src/kernel.cu @@ -1,4 +1,10 @@ #define GLM_FORCE_CUDA + +#include +#include +#include +#include + #include #include #include @@ -85,6 +91,7 @@ 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_pos2; // LOOK-2.1 - Grid parameters based on simulation parameters. // These are automatically computed for you in Boids::initSimulation @@ -92,6 +99,7 @@ int gridCellCount; int gridSideCount; float gridCellWidth; float gridInverseCellWidth; +float gridNeighborhoodDistance; glm::vec3 gridMinimum; /****************** @@ -157,7 +165,8 @@ void Boids::initSimulation(int N) { checkCUDAErrorWithLine("kernGenerateRandomPosArray failed!"); // LOOK-2.1 computing grid params - gridCellWidth = 2.0f * std::max(std::max(rule1Distance, rule2Distance), rule3Distance); + gridNeighborhoodDistance = std::max(std::max(rule1Distance, rule2Distance), rule3Distance); + gridCellWidth = 2.0f * gridNeighborhoodDistance; int halfSideCount = (int)(scene_scale / gridCellWidth) + 1; gridSideCount = 2 * halfSideCount; @@ -169,6 +178,23 @@ void Boids::initSimulation(int N) { gridMinimum.z -= halfGridWidth; // TODO-2.1 TODO-2.3 - Allocate additional buffers here. + cudaMalloc((void**)&dev_particleArrayIndices, N * sizeof(int)); + checkCUDAErrorWithLine("cudaMalloc dev_particleArrayIndices failed!"); + dev_thrust_particleArrayIndices = thrust::device_ptr(dev_particleArrayIndices); + + cudaMalloc((void**)&dev_particleGridIndices, N * sizeof(int)); + checkCUDAErrorWithLine("cudaMalloc dev_particleGridIndices failed!"); + dev_thrust_particleGridIndices = thrust::device_ptr(dev_particleGridIndices); + + cudaMalloc((void**)&dev_gridCellStartIndices, gridCellCount * sizeof(int)); + checkCUDAErrorWithLine("cudaMalloc dev_gridCellStartIndices failed!"); + + cudaMalloc((void**)&dev_gridCellEndIndices, gridCellCount * sizeof(int)); + checkCUDAErrorWithLine("cudaMalloc dev_gridCellEndIndices failed!"); + + cudaMalloc((void**)&dev_pos2, N * sizeof(glm::vec3)); + checkCUDAErrorWithLine("cudaMalloc dev_pos2 failed!"); + cudaDeviceSynchronize(); } @@ -230,10 +256,63 @@ 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) { + glm::vec3 velocityDelta = glm::vec3(0.f); + // Rule 1: boids fly towards their local perceived center of mass, which excludes themselves + glm::vec3 perceivedCenter = glm::vec3(0.f); + int rule1NumNeighbors = 0; + + for (int i = 0; i < N; ++i) + { + if (i == iSelf && glm::distance(pos[iSelf], pos[i]) < rule1Distance) + { + perceivedCenter += pos[i]; + ++rule1NumNeighbors; + } + } + + if (rule1NumNeighbors > 0) + { + perceivedCenter /= rule1NumNeighbors; + } + + velocityDelta += (perceivedCenter - pos[iSelf]) * rule1Scale; + // Rule 2: boids try to stay a distance d away from each other + glm::vec3 c = glm::vec3(0.f); + + for (int i = 0; i < N; ++i) + { + if (i != iSelf && glm::distance(pos[iSelf], pos[i]) < rule2Distance) + { + c -= (pos[i] - pos[iSelf]); + } + } + + velocityDelta += c * rule2Scale; + // Rule 3: boids try to match the speed of surrounding boids - return glm::vec3(0.0f, 0.0f, 0.0f); + glm::vec3 perceivedVelocity = glm::vec3(0.f); + int rule3NumNeighbors = 0; + + for (int i = 0; i < N; ++i) + { + if (i != iSelf && glm::distance(pos[iSelf], pos[i]) < rule3Distance) + { + perceivedVelocity += vel[i]; + ++rule3NumNeighbors; + } + } + + if (rule3NumNeighbors > 0) + { + perceivedVelocity /= rule3NumNeighbors; + } + + velocityDelta += perceivedVelocity * rule3Scale; + + // Output the combined velocity change + return velocityDelta; } /** @@ -242,9 +321,29 @@ __device__ glm::vec3 computeVelocityChange(int N, int iSelf, const glm::vec3 *po */ __global__ void kernUpdateVelocityBruteForce(int N, glm::vec3 *pos, glm::vec3 *vel1, glm::vec3 *vel2) { + int index = threadIdx.x + (blockIdx.x * blockDim.x); + + if (index >= N) + { + return; + } + // Compute a new velocity based on pos and vel1 + glm::vec3 velocityDelta = computeVelocityChange(N, index, pos, vel1); + glm::vec3 newVelocity = vel1[index] + velocityDelta; + // Clamp the speed + if (glm::length(newVelocity) > maxSpeed) + { + newVelocity = glm::normalize(newVelocity) * maxSpeed; + } + // Record the new velocity into vel2. Question: why NOT vel1? + // Answer: The new velocity depends on the previous state's velocity. + // vel1 contains the previous velocities so: + // If we were to overwrite vel1 values... + // we would no longer have the previous state's velocity + vel2[index] = newVelocity; } /** @@ -282,13 +381,27 @@ __device__ int gridIndex3Dto1D(int x, int y, int z, int gridResolution) { return x + y * gridResolution + z * gridResolution * gridResolution; } +__device__ glm::ivec3 posToGridIndex3D(const glm::vec3 pos, const glm::vec3 gridMin, const float inverseCellWidth) +{ + return glm::ivec3((pos - gridMin) * inverseCellWidth); +} + __global__ void kernComputeIndices(int N, int gridResolution, glm::vec3 gridMin, float inverseCellWidth, glm::vec3 *pos, int *indices, int *gridIndices) { // TODO-2.1 + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= N) return; + // - Label each boid with the index of its grid cell. + glm::vec3 iPos = pos[index]; + glm::ivec3 gridPos = glm::ivec3((iPos - gridMin) * inverseCellWidth); + int gridIndex = gridIndex3Dto1D(gridPos.x, gridPos.y, gridPos.z, gridResolution); + gridIndices[index] = gridIndex; + // - Set up a parallel array of integer indices as pointers to the actual // boid data in pos and vel1/vel2 + indices[index] = index; } // LOOK-2.1 Consider how this could be useful for indicating that a cell @@ -306,41 +419,234 @@ __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 = threadIdx.x + (blockIdx.x * blockDim.x); + if (0 >= index || index >= N) return; + + int prev = particleGridIndices[index - 1]; + int curr = particleGridIndices[index]; + + if (prev != curr) + { + gridCellEndIndices[prev] = index - 1; + gridCellStartIndices[curr] = index; + } + + if (index == N - 1) + { + gridCellEndIndices[curr] = index; + } +} + +__global__ void kernSortCoherentGrid(const int N, const int* particleArrayIndices, const glm::vec3* pos, const glm::vec3* vel, glm::vec3* outPos, glm::vec3* outVel) +{ + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= N) return; + + outPos[index] = pos[particleArrayIndices[index]]; + outVel[index] = vel[particleArrayIndices[index]]; } __global__ void kernUpdateVelNeighborSearchScattered( int N, int gridResolution, glm::vec3 gridMin, - float inverseCellWidth, float cellWidth, + float inverseCellWidth, float cellWidth, float neighborhoodDistance, int *gridCellStartIndices, int *gridCellEndIndices, int *particleArrayIndices, glm::vec3 *pos, glm::vec3 *vel1, glm::vec3 *vel2) { // TODO-2.1 - Update a boid's velocity using the uniform grid to reduce // the number of boids that need to be checked. + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= N) return; + // - Identify the grid cell that this particle is in + glm::vec3 iPos = pos[index]; + glm::ivec3 gridPos = posToGridIndex3D(iPos, gridMin, inverseCellWidth); + // - 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. + glm::ivec3 neighborCellStart = posToGridIndex3D(iPos - glm::vec3(neighborhoodDistance), gridMin, inverseCellWidth); + glm::ivec3 neighborCellEnd = posToGridIndex3D(iPos + glm::vec3(neighborhoodDistance), gridMin, inverseCellWidth); + + neighborCellStart = glm::max(neighborCellStart, 0); + neighborCellEnd = glm::min(neighborCellEnd, gridResolution - 1); + + int rule1NumNeighbors = 0, rule3NumNeighbors = 0; + glm::vec3 perceivedCenter = glm::vec3(0.f), c = glm::vec3(0.f), perceivedVelocity = glm::vec3(0.f); + glm::vec3 posSelf = pos[index]; + + for (int z = neighborCellStart.z; z <= neighborCellEnd.z; ++z) + { + for (int y = neighborCellStart.y; y <= neighborCellEnd.y; ++y) + { + for (int x = neighborCellStart.x; x <= neighborCellEnd.x; ++x) + { + // - For each cell, read the start/end indices in the boid pointer array. + int cellIndex = gridIndex3Dto1D(x, y, z, gridResolution); + if (cellIndex < 0 || cellIndex >= gridResolution * gridResolution * gridResolution) continue; + + int gridCellStartIndex = gridCellStartIndices[cellIndex]; + int gridCellEndIndex = gridCellEndIndices[cellIndex]; + + // - Access each boid in the cell and compute velocity change from + // the boids rules, if this boid is within the neighborhood distance. + for (int i = gridCellStartIndex; i <= gridCellEndIndex; ++i) + { + int iOther = particleArrayIndices[i]; + + if (iOther == index) continue; + glm::vec3 posOther = pos[iOther]; + float distance = glm::distance(posSelf, posOther); + + // Rule 1 accumulators + if (distance < rule1Distance) + { + perceivedCenter += posOther; + ++rule1NumNeighbors; + } + + // Rule 2 accumulators + if (distance < rule2Distance) + { + c -= (posOther - posSelf); + } + + // Rule 3 accumulators + if (distance < rule3Distance) + { + perceivedVelocity += vel1[iOther]; + ++rule3NumNeighbors; + } + } + } + } + } + + glm::vec3 velocityDelta = glm::vec3(0.f); + + // Rule 1 velocity delta + if (rule1NumNeighbors > 0) + { + perceivedCenter /= rule1NumNeighbors; + velocityDelta += (perceivedCenter - pos[index]) * rule1Scale; + } + + // Rule 2 velocity delta + velocityDelta += c * rule2Scale; + + // Rule 3 velocity delta + if (rule3NumNeighbors > 0) + { + perceivedVelocity /= rule3NumNeighbors; + velocityDelta += perceivedVelocity * rule3Scale; + } + // - Clamp the speed change before putting the new speed in vel2 + glm::vec3 newVel = vel1[index] + velocityDelta; + + if (glm::length(newVel) > maxSpeed) + { + newVel = glm::normalize(newVel) * maxSpeed; + } + + vel2[index] = newVel; } __global__ void kernUpdateVelNeighborSearchCoherent( int N, int gridResolution, glm::vec3 gridMin, - float inverseCellWidth, float cellWidth, + float inverseCellWidth, float cellWidth, float neighborhoodDistance, int *gridCellStartIndices, int *gridCellEndIndices, glm::vec3 *pos, glm::vec3 *vel1, glm::vec3 *vel2) { // TODO-2.3 - This should be very similar to kernUpdateVelNeighborSearchScattered, // except with one less level of indirection. // This should expect gridCellStartIndices and gridCellEndIndices to refer // directly to pos and vel1. + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= N) return; + // - Identify the grid cell that this particle is in + glm::vec3 iPos = pos[index]; + glm::ivec3 gridPos = posToGridIndex3D(iPos, gridMin, inverseCellWidth); + // - 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. + glm::ivec3 neighborCellStart = posToGridIndex3D(iPos - glm::vec3(neighborhoodDistance), gridMin, inverseCellWidth); + glm::ivec3 neighborCellEnd = posToGridIndex3D(iPos + glm::vec3(neighborhoodDistance), gridMin, inverseCellWidth); + + neighborCellStart = glm::max(neighborCellStart, 0); + neighborCellEnd = glm::min(neighborCellEnd, gridResolution - 1); + + int rule1NumNeighbors = 0, rule3NumNeighbors = 0; + glm::vec3 perceivedCenter = glm::vec3(0.f), c = glm::vec3(0.f), perceivedVelocity = glm::vec3(0.f); + glm::vec3 posSelf = pos[index]; + + for (int z = neighborCellStart.z; z <= neighborCellEnd.z; ++z) + { + for (int y = neighborCellStart.y; y <= neighborCellEnd.y; ++y) + { + for (int x = neighborCellStart.x; x <= neighborCellEnd.x; ++x) + { + // - 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. + int cellIndex = gridIndex3Dto1D(x, y, z, gridResolution); + + int gridCellStartIndex = gridCellStartIndices[cellIndex]; + int gridCellEndIndex = gridCellEndIndices[cellIndex]; + + // - Access each boid in the cell and compute velocity change from + // the boids rules, if this boid is within the neighborhood distance. + for (int i = gridCellStartIndex; i <= gridCellEndIndex; ++i) + { + if (i == index) continue; + glm::vec3 posOther = pos[i]; + float distance = glm::distance(posSelf, posOther); + + // Rule 1 + if (distance < rule1Distance) + { + perceivedCenter += posOther; + ++rule1NumNeighbors; + } + + // Rule 2 + if (distance < rule2Distance) + { + c -= (posOther - posSelf); + } + + // Rule 3 + if (distance < rule3Distance) + { + perceivedVelocity += vel1[i]; + ++rule3NumNeighbors; + } + } + } + } + } + + glm::vec3 velocityDelta = glm::vec3(0.f); + + if (rule1NumNeighbors > 0) + { + perceivedCenter /= rule1NumNeighbors; + velocityDelta += (perceivedCenter - pos[index]) * rule1Scale; + } + + velocityDelta += c * rule2Scale; + + if (rule3NumNeighbors > 0) + { + perceivedVelocity /= rule3NumNeighbors; + velocityDelta += perceivedVelocity * rule3Scale; + } + // - Clamp the speed change before putting the new speed in vel2 + glm::vec3 newVel = vel1[index] + velocityDelta; + + if (glm::length(newVel) > maxSpeed) + { + newVel = glm::normalize(newVel) * maxSpeed; + } + + vel2[index] = newVel; } /** @@ -348,40 +654,98 @@ __global__ void kernUpdateVelNeighborSearchCoherent( */ void Boids::stepSimulationNaive(float dt) { // TODO-1.2 - use the kernels you wrote to step the simulation forward in time. + dim3 fullBlocksPerGrid((numObjects + blockSize - 1) / blockSize); + kernUpdateVelocityBruteForce<<>>(numObjects, dev_pos, dev_vel1, dev_vel2); + kernUpdatePos<<>>(numObjects, dt, dev_pos, dev_vel2); // TODO-1.2 ping-pong the velocity buffers + glm::vec3 *tmp_vel = dev_vel1; + dev_vel1 = dev_vel2; + dev_vel2 = tmp_vel; } void Boids::stepSimulationScatteredGrid(float dt) { // TODO-2.1 // Uniform Grid Neighbor search using Thrust sort. // In Parallel: + dim3 fullBoidBlocksPerGrid((numObjects + blockSize - 1) / blockSize); + dim3 fullCellBlocksPerGrid((gridCellCount + blockSize - 1) / blockSize); + // - 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); + // - 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, 0); + kernResetIntBuffer<<>>(gridCellCount, dev_gridCellEndIndices, -1); + kernIdentifyCellStartEnd<<>>(numObjects, dev_particleGridIndices, dev_gridCellStartIndices, dev_gridCellEndIndices); + // - Perform velocity updates using neighbor search + kernUpdateVelNeighborSearchScattered<<>>( + numObjects, + gridSideCount, gridMinimum, gridInverseCellWidth, gridCellWidth, gridNeighborhoodDistance, + dev_gridCellStartIndices, dev_gridCellEndIndices, dev_particleArrayIndices, + dev_pos, dev_vel1, dev_vel2); + // - Update positions + kernUpdatePos<<>>(numObjects, dt, dev_pos, dev_vel2); + // - Ping-pong buffers as needed + glm::vec3 *tmp_vel = dev_vel1; + dev_vel1 = dev_vel2; + dev_vel2 = tmp_vel; } void Boids::stepSimulationCoherentGrid(float dt) { // TODO-2.3 - start by copying Boids::stepSimulationNaiveGrid // Uniform Grid Neighbor search using Thrust sort on cell-coherent data. // In Parallel: + dim3 fullBoidBlocksPerGrid((numObjects + blockSize - 1) / blockSize); + dim3 fullCellBlocksPerGrid((gridCellCount + blockSize - 1) / blockSize); + // - 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); + // - 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, 0); + kernResetIntBuffer<<>>(gridCellCount, dev_gridCellEndIndices, -1); + kernIdentifyCellStartEnd<<>>(numObjects, dev_particleGridIndices, dev_gridCellStartIndices, dev_gridCellEndIndices); + // - BIG DIFFERENCE: use the rearranged array index buffer to reshuffle all // the particle data in the simulation array. // CONSIDER WHAT ADDITIONAL BUFFERS YOU NEED + kernSortCoherentGrid<<>>(numObjects, dev_particleArrayIndices, dev_pos, dev_vel1, dev_pos2, dev_vel2); + // - Perform velocity updates using neighbor search + kernUpdateVelNeighborSearchCoherent<<>>( + numObjects, + gridSideCount, gridMinimum, gridInverseCellWidth, gridCellWidth, gridNeighborhoodDistance, + dev_gridCellStartIndices, dev_gridCellEndIndices, + dev_pos2, dev_vel2, dev_vel1); + // - Update positions + kernUpdatePos<<>>(numObjects, dt, dev_pos2, dev_vel2); + // - Ping-pong buffers as needed. THIS MAY BE DIFFERENT FROM BEFORE. + glm::vec3* tmp_pos = dev_pos; + dev_pos = dev_pos2; + dev_pos2 = tmp_pos; } void Boids::endSimulation() { @@ -390,11 +754,18 @@ 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_pos2); } void Boids::unitTest() { + return; // LOOK-1.2 Feel free to write additional tests here. - // test unstable sort int *dev_intKeys; int *dev_intValues; @@ -449,9 +820,35 @@ void Boids::unitTest() { std::cout << " value: " << intValues[i] << std::endl; } + // Test identify cell start and end + std::unique_ptr startIndices{ new int[N] }; + std::unique_ptr endIndices{ new int[N] }; + int *dev_startIndices, *dev_endIndices; + + cudaMalloc((void**)&dev_startIndices, N * sizeof(int)); + checkCUDAErrorWithLine("cudaMalloc dev_startIndices failed!"); + + cudaMalloc((void**)&dev_endIndices, N * sizeof(int)); + checkCUDAErrorWithLine("cudaMalloc dev_endIndices failed!"); + + kernResetIntBuffer<<<1, N>>>(N, dev_startIndices, 0); + kernResetIntBuffer<<<1, N>>>(N, dev_endIndices, -1); + kernIdentifyCellStartEnd<<<1, N>>>(N, dev_intKeys, dev_startIndices, dev_endIndices); + + cudaMemcpy(startIndices.get(), dev_startIndices, sizeof(int) * N, cudaMemcpyDeviceToHost); + cudaMemcpy(endIndices.get(), dev_endIndices, sizeof(int) * N, cudaMemcpyDeviceToHost); + + std::cout << "start and end indices: " << std::endl; + for (int i = 0; i < N; i++) { + std::cout << " key: " << i; + std::cout << " value: [" << startIndices[i] << ", " << endIndices[i] << "]" << std::endl; + } + // cleanup cudaFree(dev_intKeys); cudaFree(dev_intValues); + cudaFree(dev_startIndices); + cudaFree(dev_endIndices); checkCUDAErrorWithLine("cudaFree failed!"); return; } diff --git a/src/kernel.h b/src/kernel.h index ab8b9ab..0ca6160 100644 --- a/src/kernel.h +++ b/src/kernel.h @@ -1,10 +1,6 @@ #pragma once #include -#include -#include -#include -#include #include #include diff --git a/src/main.cpp b/src/main.cpp index fe657ed..bc00324 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -17,8 +17,8 @@ // 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;