+ 
+ 
+ 
+
+
+
+
+
+
+
+
+
+
+
+### Questions
+* For each implementation, how does changing the number of boids affect performance? Why do you think this is?
+
+
+ For each implementation, the increasing number of boids will result in decreasing performance because there are simply more calculation as number of boids increases. The naive method will loop through all the boids to find its neighbors. This method is the most affected by the number of boids and it has steeper decrease as the number of boids increases. For the two grid methods, they save a lot of calculation by preproccessing the space into grid. They both have a smoothier decrease between 5000 and 10000 boids. I think the memory read overhead for 5000 and 10000 can still cover for all the neighbor-finding iteration. As the number of boid increases, this benefit start to vanish and still result in steep decrease in performance.
+
+
+* For each implementation, how does changing the block count and block size affect performance? Why do you think this is?
+
+ For all 3 implementations, changning the block size only affect performance a little bit, looking at the chart. However, there is a slightly noticeable decrease when the block size drop to 32. I think this is when we are under-utilizing the GPU resources and more adding overhead of launching more thread per block.
+
+* 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?
+
+ In general, the coherent uniform grid perform better than the scattered uniform grid method if you look at the chart. With or without visualization, the coherent uniform grid method has higher FPS than the scattered uniform grid under 128 block size. This is expected becasue even though we spend some performance to rearrange the position and velocity array, reading in a contiguous memory will be much faster than reading scattered memeory.
+
+* 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!
+
+ For my implementation, checking 27 neighboring cells does result in worse performance. I checked with 50000 boids using 128 block size. With 8 ceslls, it's hitting around 980 FPS, while 27 cells can only hit around 675 FPS. I also checked with smaller and larger number of boids and got the same conclusion. Other than the reason that there are more cells to check, I think fewer cells in conherent uniform grid method also have the advantage that closer the neighboring cell to the current boid, their memory location are actually closer, so it will cost less time to fetch those data than gathering data for outer grid cells which will locate in further memory location. Checking fewer cells not only results in less iteration but also takes less time to read memory.
diff --git a/images/blocksize.png b/images/blocksize.png
new file mode 100644
index 0000000..345c11e
Binary files /dev/null and b/images/blocksize.png differ
diff --git a/images/demo-10000.gif b/images/demo-10000.gif
new file mode 100644
index 0000000..5176de5
Binary files /dev/null and b/images/demo-10000.gif differ
diff --git a/images/demo-100000.gif b/images/demo-100000.gif
new file mode 100644
index 0000000..92cebc1
Binary files /dev/null and b/images/demo-100000.gif differ
diff --git a/images/demo-5000-1.gif b/images/demo-5000-1.gif
new file mode 100644
index 0000000..858ca4f
Binary files /dev/null and b/images/demo-5000-1.gif differ
diff --git a/images/demo-5000.gif b/images/demo-5000.gif
new file mode 100644
index 0000000..9eb18a0
Binary files /dev/null and b/images/demo-5000.gif differ
diff --git a/images/demo-50000.gif b/images/demo-50000.gif
new file mode 100644
index 0000000..acac4e7
Binary files /dev/null and b/images/demo-50000.gif differ
diff --git a/images/vis_off.png b/images/vis_off.png
new file mode 100644
index 0000000..9c29b28
Binary files /dev/null and b/images/vis_off.png differ
diff --git a/images/vis_on.png b/images/vis_on.png
new file mode 100644
index 0000000..1fe2b6c
Binary files /dev/null and b/images/vis_on.png differ
diff --git a/src/kernel.cu b/src/kernel.cu
index 74dffcb..17a955e 100644
--- a/src/kernel.cu
+++ b/src/kernel.cu
@@ -85,6 +85,8 @@ 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_coherent_vel;
+glm::vec3* dev_coherent_pos;
// LOOK-2.1 - Grid parameters based on simulation parameters.
// These are automatically computed for you in Boids::initSimulation
@@ -169,6 +171,24 @@ 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!");
+
+ 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_coherent_vel, N * sizeof(glm::vec3));
+ checkCUDAErrorWithLine("cudaMalloc dev_coherent_vel failed!");
+
+ cudaMalloc((void**)&dev_coherent_pos, N * sizeof(glm::vec3));
+ checkCUDAErrorWithLine("cudaMalloc dev_coherent_pos failed!");
+
cudaDeviceSynchronize();
}
@@ -229,11 +249,83 @@ 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 iSelf, const glm::vec3* pos, const glm::vec3* vel, const int* all_neighbor, int neighbor_count) {
+ glm::vec3 res = vel[iSelf];
+ glm::vec3 iPos = pos[iSelf];
+ glm::vec3 rule1Sum = glm::vec3(0.0f);
+ glm::vec3 rule2Sum = glm::vec3(0.0f);
+ glm::vec3 rule3Sum = glm::vec3(0.0f);
+ int numOfNeighor = 0;
+
+ for (int i = 0; i < neighbor_count; i++) {
+
+ float distance = glm::distance(iPos, pos[all_neighbor[i]]);
+ if (distance < rule1Distance) {
+ rule1Sum += pos[all_neighbor[i]];
+ numOfNeighor++;
+
+ if (distance < rule2Distance) {
+ rule2Sum -= (pos[all_neighbor[i]] - iPos);
+ }
+ if (distance < rule3Distance) {
+ rule3Sum += vel[all_neighbor[i]];
+ }
+ }
+ }
+
+ if (numOfNeighor > 0) {
+ rule1Sum /= numOfNeighor;
+ res += (rule1Sum - iPos) * rule1Scale;
+
+ rule3Sum /= numOfNeighor;
+ res += rule3Sum * rule3Scale;
+ }
+
+ res += rule2Sum * rule2Scale;
+
+ return res;
+}
+
__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 res = vel[iSelf];
+ glm::vec3 iPos = pos[iSelf];
+ glm::vec3 rule1Sum = glm::vec3(0.0f);
+ glm::vec3 rule2Sum = glm::vec3(0.0f);
+ glm::vec3 rule3Sum = glm::vec3(0.0f);
+ int numOfNeighor = 0;
+
+ for (int i = 0; i < N; i++) {
+ if (i == iSelf) continue;
+
+ float distance = glm::distance(iPos, pos[i]);
+ if (distance < rule1Distance) {
+ rule1Sum += pos[i];
+ numOfNeighor++;
+
+ if (distance < rule2Distance) {
+ rule2Sum -= (pos[i] - iPos);
+ }
+ if (distance < rule3Distance) {
+ rule3Sum += vel[i];
+ }
+ }
+ }
+
+ if (numOfNeighor > 0) {
+ rule1Sum /= numOfNeighor;
+ res += (rule1Sum - iPos) * rule1Scale;
+
+ rule3Sum /= numOfNeighor;
+ res += rule3Sum * rule3Scale;
+ }
+
+ res += rule2Sum * rule2Scale;
+
+ return res;
}
/**
@@ -245,6 +337,18 @@ __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 newVel = computeVelocityChange(N, index, pos, vel1);
+
+ float len = glm::length(newVel);
+ if (len > maxSpeed) {
+ newVel = (newVel / len) * maxSpeed;
+ }
+
+ vel2[index] = newVel;
}
/**
@@ -289,6 +393,13 @@ __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) {
+ glm::vec3 boid_pos = pos[index];
+ glm::vec3 boid_in_grid = floor((boid_pos - gridMin) * inverseCellWidth);
+ gridIndices[index] = gridIndex3Dto1D(boid_in_grid.x, boid_in_grid.y, boid_in_grid.z, gridResolution);
+ indices[index] = index;
+ }
}
// LOOK-2.1 Consider how this could be useful for indicating that a cell
@@ -306,6 +417,17 @@ __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) {
+ int cur_grid_idx = particleGridIndices[index];
+ if (index == 0 || cur_grid_idx != particleGridIndices[index - 1]) {
+ gridCellStartIndices[cur_grid_idx] = index;
+ }
+
+ if (index == N - 1 || cur_grid_idx != particleGridIndices[index + 1]) {
+ gridCellEndIndices[cur_grid_idx] = index;
+ }
+ }
}
__global__ void kernUpdateVelNeighborSearchScattered(
@@ -322,6 +444,77 @@ __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 i = threadIdx.x + (blockIdx.x * blockDim.x);
+ if (i >= N) {
+ return;
+ }
+
+ glm::vec3 res = vel1[i];
+ glm::vec3 iPos = pos[i];
+ glm::vec3 rule1Sum = glm::vec3(0.0f);
+ glm::vec3 rule2Sum = glm::vec3(0.0f);
+ glm::vec3 rule3Sum = glm::vec3(0.0f);
+ int numOfNeighor = 0;
+
+
+ //find its neighbors's grid index
+ for (int x = -1; x <= 1; x++) {
+ for (int y = -1; y <= 1; y++) {
+ for (int z = -1; z <= 1; z++) {
+ glm::vec3 neighbor_range_pos = iPos + glm::vec3(x * cellWidth, y * cellWidth, z * cellWidth);
+ glm::vec3 boid_in_grid = floor((neighbor_range_pos - gridMin) * inverseCellWidth);
+ int neighbor_grid_idx = gridIndex3Dto1D(static_cast