diff --git a/README.md b/README.md index 98dd9a8..97abefc 100644 --- a/README.md +++ b/README.md @@ -1,10 +1,82 @@ **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) +* LINSHEN XIAO +* Tested on: Windows 10, Intel(R) Core(TM) i7-6700HQ CPU @ 2.60GHz, 16.0GB, NVIDIA GeForce GTX 970M (Personal computer) -### (TODO: Your README) +## Screenshot -Include screenshots, analysis, etc. (Remember, this is public, so don't put -anything here that you don't want to share with the world.) +A screenshot of the boids: + +![](images/boids.gif) + +## Performance analysis + +* Framerate change with increasing # of boids for naive, scattered uniform grid, and coherent uniform grid (with and without visualization) + +![](images/Withvisualize.png) + +Boids | Navie | Scattered uniform grid | Coherent uniform grid +:---|:---:|:---:|:---: +5000 | 156.2 | 397 | 401.8 +10000 | 55.7 | 481.6 | 455.9 +50000 | 2.8 | 278.6 | 341.1 +100000 | 0.7 | 149.5 | 185.2 +500000 | ? | 21.9 | 35.3 + +![](images/Withoutvisualize.png) + +Boids | Navie | Scattered uniform grid | Coherent uniform grid +:---|:---:|:---:|:---: +5000 | 210.5 | 810.4 | 728.0 +10000 | 63.4 | 1174.7 | 1231.8 +50000 | 2.8 | 450.2 | 553.4 +100000 | 0.7 | 205.6 | 294.9 +500000 | ? | 22.3| 40.2 + +As the number of boids grows, the FPS may always decrease(with some small exception), that's because more data needs to be computed if the number of boids increase. + +Improvement: Coherent uniform grid > Scattered uniform grid > Naive + +Framerate change with increasing block size + +![](images/Blocksize.png) + +Blocksize | Coherent uniform grid +:---|:---: +16 | 228 +32 | 309 +64 | 318.2 +128 | 307.6 +256 | 326.4 +512 | 316.9 + +The change of block size does not affect the performance dramatically, but we can still see that block size 32 is still better than block size 16. + +## Questions + +* For each implementation, how does changing the number of boids affect performance? Why do you think this is? + +As the number of boids grows, the FPS will decrease, that's because more data needs to be computed if the number of boids increase. + +* For each implementation, how does changing the block count and block size affect performance? Why do you think this is? + +The change of block count and block size does not affect the performance dramatically, but it has some improvement before 32. + +* 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. We can see the great improvement when the number of boids > 50000. Coherent uniform grid sort the array to skip a searching step to improve the performance. + +* Did changing cell width and checking 27 vs 8 neighboring cells affect performance? Why or why not? + +![](images/27vs8.png) + +Boids | 27 Cells | 8 Cells +:---|:---:|:---: +5000 | 412.8 | 401.8 +10000 | 482.6 | 455.9 +50000 | 450 | 341.1 +100000 | 303.2 | 185.2 +500000 | 97.8 | 35.3 + +Yes. When the boids' amount is huge, though 27 cells needs to iterates more neighbors, the number of boids needs to be interated is less, so a large amount of time of computing distance is saved. \ No newline at end of file diff --git a/images/27vs8.png b/images/27vs8.png new file mode 100644 index 0000000..f4696e6 Binary files /dev/null and b/images/27vs8.png differ diff --git a/images/Blocksize.png b/images/Blocksize.png new file mode 100644 index 0000000..13022d0 Binary files /dev/null and b/images/Blocksize.png differ diff --git a/images/Withoutvisualize.png b/images/Withoutvisualize.png new file mode 100644 index 0000000..f8a3971 Binary files /dev/null and b/images/Withoutvisualize.png differ diff --git a/images/Withvisualize.png b/images/Withvisualize.png new file mode 100644 index 0000000..7955467 Binary files /dev/null and b/images/Withvisualize.png differ diff --git a/images/boids.gif b/images/boids.gif new file mode 100644 index 0000000..38e196b Binary files /dev/null and b/images/boids.gif differ diff --git a/src/kernel.cu b/src/kernel.cu index aaf0fbf..346a8ec 100644 --- a/src/kernel.cu +++ b/src/kernel.cu @@ -85,6 +85,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 @@ -156,6 +157,7 @@ void Boids::initSimulation(int N) { dev_pos, scene_scale); checkCUDAErrorWithLine("kernGenerateRandomPosArray failed!"); + //27:1.0f/8:2.0f // LOOK-2.1 computing grid params gridCellWidth = 2.0f * std::max(std::max(rule1Distance, rule2Distance), rule3Distance); int halfSideCount = (int)(scene_scale / gridCellWidth) + 1; @@ -169,6 +171,25 @@ 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!"); + + cudaMalloc((void**)&dev_particleGridIndices, N * sizeof(int)); + checkCUDAErrorWithLine("cudaMalloc dev_particleGridIndices failed!"); + + dev_thrust_particleArrayIndices = thrust::device_ptr(dev_particleArrayIndices); + 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!"); + cudaThreadSynchronize(); } @@ -230,10 +251,52 @@ 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 rule1vel(0.f); + glm::vec3 rule2vel(0.f); + glm::vec3 rule3vel(0.f); + glm::vec3 perceived_center(0.f); + glm::vec3 c(0.f); + glm::vec3 perceived_velocity(0.f); + int rule1num = 0; + int rule3num = 0; + for (int i = 0; i < N; i++) + { + // Rule 1: boids fly towards their local perceived center of mass, which excludes themselves + if (i != iSelf) + { + float distance = glm::length(pos[i] - pos[iSelf]); + if (distance < rule1Distance) + { + perceived_center += pos[i]; + rule1num++; + } + // Rule 2: boids try to stay a distance d away from each other + if (distance < rule2Distance) + { + c -= (pos[i] - pos[iSelf]); + } + // Rule 3: boids try to match the speed of surrounding boids + if (distance < rule3Distance) + { + perceived_velocity += vel[i]; + rule3num++; + } + } + } + if (rule1num > 0) + { + perceived_center /= float(rule1num); + rule1vel = (perceived_center - pos[iSelf])*rule1Scale; + } + + rule2vel = c * rule2Scale; + + if (rule3num > 0) + { + perceived_velocity /= float(rule3num); + rule3vel = perceived_velocity * rule3Scale; + } + return rule1vel + rule2vel + rule3vel; } /** @@ -242,9 +305,19 @@ __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 newvel = vel1[index] + computeVelocityChange(N, index, pos, vel1); // Clamp the speed + if (glm::length(newvel) > maxSpeed) + { + newvel = glm::normalize(newvel)*maxSpeed; + } // Record the new velocity into vel2. Question: why NOT vel1? + vel2[index] = newvel; } /** @@ -279,7 +352,14 @@ __global__ void kernUpdatePos(int N, float dt, glm::vec3 *pos, glm::vec3 *vel) { // for(y) // for(z)? Or some other order? __device__ int gridIndex3Dto1D(int x, int y, int z, int gridResolution) { - return x + y * gridResolution + z * gridResolution * gridResolution; + if (x<0 || x>gridResolution || x<0 || x>gridResolution || x<0 || x>gridResolution) + { + return -1; + } + else + { + return x + y * gridResolution + z * gridResolution * gridResolution; + } } __global__ void kernComputeIndices(int N, int gridResolution, @@ -289,6 +369,14 @@ __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 = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= N) + { + return; + } + glm::vec3 gridIndices3D = (pos[index] - gridMin) * inverseCellWidth; + gridIndices[index] = gridIndex3Dto1D((int)gridIndices3D.x, (int)gridIndices3D.y, (int)gridIndices3D.z, gridResolution); + indices[index] = index; } // LOOK-2.1 Consider how this could be useful for indicating that a cell @@ -306,6 +394,39 @@ __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 (index >= N) + { + return; + } + int gridindex = particleGridIndices[index]; + if (index == 0) + { + gridCellStartIndices[gridindex] = index; + return; + } + if (index == N - 1) + { + gridCellEndIndices[gridindex] = index; + } + int previousgridindex = particleGridIndices[index - 1]; + if (gridindex != previousgridindex) + { + gridCellStartIndices[gridindex] = index; + gridCellEndIndices[previousgridindex] = index - 1; + } +} + +/** +* use the rearranged array index buffer to reshuffle all +* the particle data in the simulation array. +*/ +__global__ void kernReshuffle(int N, int *particleGridIndices, glm::vec3 *ori, glm::vec3 *tg) { + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= N) { + return; + } + tg[index] = ori[particleGridIndices[index]]; } __global__ void kernUpdateVelNeighborSearchScattered( @@ -316,12 +437,96 @@ __global__ void kernUpdateVelNeighborSearchScattered( 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 + int boidindex = particleArrayIndices[index]; + glm::vec3 gridIndices3D = (pos[boidindex] - gridMin) * inverseCellWidth; // - Identify which cells may contain neighbors. This isn't always 8. + int neighborcell[8]; + int indexx = int(gridIndices3D.x); + int indexy = int(gridIndices3D.y); + int indexz = int(gridIndices3D.z); + int offsetx = (gridIndices3D.x - indexx) >= 0.5f ? 1 : -1; + int offsety = (gridIndices3D.y - indexy) >= 0.5f ? 1 : -1; + int offsetz = (gridIndices3D.z - indexz) >= 0.5f ? 1 : -1; + int count = 0; + for (int i = 0; i < 2; i++) + { + for (int j = 0; j < 2; j++) + { + for (int k = 0; k < 2; k++) + { + neighborcell[count] = gridIndex3Dto1D(indexx + i*offsetx, indexy + j*offsety, indexz + k*offsetz, gridResolution); + count++; + } + } + } + glm::vec3 rule1vel(0.f); + glm::vec3 rule2vel(0.f); + glm::vec3 rule3vel(0.f); + glm::vec3 perceived_center(0.f); + glm::vec3 c(0.f); + glm::vec3 perceived_velocity(0.f); + int rule1num = 0; + int rule3num = 0; // - 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. + for (int i = 0; i < 8; i++) + { + if (neighborcell[i] == -1) + { + continue; + } + for (int j = gridCellStartIndices[neighborcell[i]]; j < gridCellEndIndices[neighborcell[i]]; j++) + { + int k = particleArrayIndices[j]; + if (k != boidindex) + { + float distance = glm::length(pos[k] - pos[boidindex]); + // Rule 1: boids fly towards their local perceived center of mass, which excludes themselves + if (distance < rule1Distance) + { + perceived_center += pos[k]; + rule1num++; + } + // Rule 2: boids try to stay a distance d away from each other + if (distance < rule2Distance) + { + c -= (pos[k] - pos[boidindex]); + } + // Rule 3: boids try to match the speed of surrounding boids + if (distance < rule3Distance) + { + perceived_velocity += vel1[k]; + rule3num++; + } + } + } + } + if (rule1num > 0) + { + perceived_center /= float(rule1num); + rule1vel = (perceived_center - pos[boidindex])*rule1Scale; + } + rule2vel = c * rule2Scale; + if (rule3num > 0) + { + perceived_velocity /= float(rule3num); + rule3vel = perceived_velocity * rule3Scale; + } + + glm::vec3 newvel = vel1[boidindex] + rule1vel + rule2vel + rule3vel; // - Clamp the speed change before putting the new speed in vel2 + if (glm::length(newvel) > maxSpeed) + { + newvel = glm::normalize(newvel)*maxSpeed; + } + vel2[boidindex] = newvel; } __global__ void kernUpdateVelNeighborSearchCoherent( @@ -333,14 +538,96 @@ __global__ void kernUpdateVelNeighborSearchCoherent( // 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 gridIndices3D = (pos[index] - gridMin) * inverseCellWidth; // - Identify which cells may contain neighbors. This isn't always 8. + int neighborcell[8]; + int indexx = int(gridIndices3D.x); + int indexy = int(gridIndices3D.y); + int indexz = int(gridIndices3D.z); + int offsetx = (gridIndices3D.x - indexx) >= 0.5f ? 1 : -1; + int offsety = (gridIndices3D.y - indexy) >= 0.5f ? 1 : -1; + int offsetz = (gridIndices3D.z - indexz) >= 0.5f ? 1 : -1; + int count = 0; + for (int i = 0; i < 2; i++) + { + for (int j = 0; j < 2; j++) + { + for (int k = 0; k < 2; k++) + { + neighborcell[count] = gridIndex3Dto1D(indexx + i*offsetx, indexy + j*offsety, indexz + k*offsetz, gridResolution); + count++; + } + } + } + glm::vec3 rule1vel(0.f); + glm::vec3 rule2vel(0.f); + glm::vec3 rule3vel(0.f); + glm::vec3 perceived_center(0.f); + glm::vec3 c(0.f); + glm::vec3 perceived_velocity(0.f); + int rule1num = 0; + int rule3num = 0; // - 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. + for (int i = 0; i < 8; i++) + { + if (neighborcell[i] == -1) + { + continue; + } + for (int j = gridCellStartIndices[neighborcell[i]]; j < gridCellEndIndices[neighborcell[i]]; j++) + { + if (j != index) + { + float distance = glm::length(pos[j] - pos[index]); + // Rule 1: boids fly towards their local perceived center of mass, which excludes themselves + if (distance < rule1Distance) + { + perceived_center += pos[j]; + rule1num++; + } + // Rule 2: boids try to stay a distance d away from each other + if (distance < rule2Distance) + { + c -= (pos[j] - pos[index]); + } + // Rule 3: boids try to match the speed of surrounding boids + if (distance < rule3Distance) + { + perceived_velocity += vel1[j]; + rule3num++; + } + } + } + } + if (rule1num > 0) + { + perceived_center /= float(rule1num); + rule1vel = (perceived_center - pos[index])*rule1Scale; + } + rule2vel = c * rule2Scale; + if (rule3num > 0) + { + perceived_velocity /= float(rule3num); + rule3vel = perceived_velocity * rule3Scale; + } + + glm::vec3 newvel = vel1[index] + rule1vel + rule2vel + rule3vel; // - Clamp the speed change before putting the new speed in vel2 + if (glm::length(newvel) > maxSpeed) + { + newvel = glm::normalize(newvel)*maxSpeed; + } + vel2[index] = newvel; } /** @@ -348,40 +635,60 @@ __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_vel1); // TODO-1.2 ping-pong the velocity buffers + std::swap(dev_vel1, dev_vel2); } void Boids::stepSimulationScatteredGrid(float dt) { // TODO-2.1 + dim3 fullBlocksPerGrid((numObjects + blockSize - 1) / blockSize); // Uniform Grid Neighbor search using Thrust sort. // 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); // - 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 + kernIdentifyCellStartEnd << > > (numObjects, dev_particleGridIndices, dev_gridCellStartIndices, dev_gridCellEndIndices); // - Perform velocity updates using neighbor search + kernUpdateVelNeighborSearchScattered << > > (numObjects, gridSideCount, gridMinimum, gridInverseCellWidth, gridCellWidth, 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 + std::swap(dev_vel1, dev_vel2); } void Boids::stepSimulationCoherentGrid(float dt) { // TODO-2.3 - start by copying Boids::stepSimulationNaiveGrid + dim3 fullBlocksPerGrid((numObjects + blockSize - 1) / blockSize); // Uniform Grid Neighbor search using Thrust sort on cell-coherent data. // 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); // - 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 + 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 + kernReshuffle << > >(numObjects, dev_particleArrayIndices, dev_pos, dev_pos2); + kernReshuffle << > >(numObjects, dev_particleArrayIndices, dev_vel2, dev_vel1); // - Perform velocity updates using neighbor search + kernUpdateVelNeighborSearchCoherent << > > (numObjects, gridSideCount, gridMinimum, gridInverseCellWidth, gridCellWidth, dev_gridCellStartIndices, dev_gridCellEndIndices, dev_pos2, dev_vel1, dev_vel2); // - Update positions + kernUpdatePos << > >(numObjects, dt, dev_pos2, dev_vel2); // - Ping-pong buffers as needed. THIS MAY BE DIFFERENT FROM BEFORE. + std::swap(dev_pos, dev_pos2); } void Boids::endSimulation() { @@ -390,6 +697,11 @@ 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() { diff --git a/src/main.cpp b/src/main.cpp index a29471d..a7c2c99 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -14,8 +14,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;