From faa6e658f72bb3d4d20fd8f8ef52d0ea53688a88 Mon Sep 17 00:00:00 2001 From: HAN YAN Date: Wed, 9 Sep 2020 13:33:00 -0400 Subject: [PATCH 01/17] init part1 --- src/kernel.cu | 67 +++++++++++++++++++++++++++++++++++++++++++++------ 1 file changed, 60 insertions(+), 7 deletions(-) diff --git a/src/kernel.cu b/src/kernel.cu index 74dffcb..5fc5e13 100644 --- a/src/kernel.cu +++ b/src/kernel.cu @@ -230,10 +230,42 @@ 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 perceived_center = glm::vec3(0.0f, 0.0f, 0.0f); + glm::vec3 c = glm::vec3(0.0f, 0.0f, 0.0f); + glm::vec3 perceived_velocity = glm::vec3(0.0f, 0.0f, 0.0f); + + for (int i = 0; i < N; i++) { + if (i == iSelf) { + continue; + } + float dist = glm::distance(pos[i], pos[iSelf]); + // Rule 1: boids fly towards their local perceived center of mass, which excludes themselves + int n1 = 0; + if (dist < rule1Distance) { + perceived_center += pos[i]; + n1++; + } + perceived_center /= n1; + + // Rule 2: boids try to stay a distance d away from each other + if (dist < rule2Distance) { + c -= (pos[i] - pos[iSelf]); + } + + // Rule 3: boids try to match the speed of surrounding boids + int n3 = 0; + if (dist < rule3Distance) { + perceived_velocity += vel[i]; + n3++; + } + perceived_velocity /= n3; + } + + glm::vec3 new_vel = vel[iSelf] + (perceived_center - pos[iSelf]) * rule1Scale + + c * rule2Scale + perceived_velocity * rule3Scale; + + return new_vel; } /** @@ -242,9 +274,24 @@ __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) { - // Compute a new velocity based on pos and vel1 - // Clamp the speed - // Record the new velocity into vel2. Question: why NOT vel1? + // Compute a new velocity based on pos and vel1 + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= N) { + return; + } + glm::vec3 new_vel = computeVelocityChange(N, index, pos, vel1); + + // Clamp the speed + new_vel.x = new_vel.x > maxSpeed ? maxSpeed : new_vel.x; + new_vel.y = new_vel.y > maxSpeed ? maxSpeed : new_vel.y; + new_vel.z = new_vel.z > maxSpeed ? maxSpeed : new_vel.z; + + new_vel.x = new_vel.x < -maxSpeed ? -maxSpeed : new_vel.x; + new_vel.y = new_vel.y < -maxSpeed ? -maxSpeed : new_vel.y; + new_vel.z = new_vel.z < -maxSpeed ? -maxSpeed : new_vel.z; + + // Record the new velocity into vel2. Question: why NOT vel1? + vel2[index] = new_vel; } /** @@ -347,8 +394,14 @@ __global__ void kernUpdateVelNeighborSearchCoherent( * Step the entire N-body simulation by `dt` seconds. */ void Boids::stepSimulationNaive(float dt) { + dim3 fullBlocksPerGrid((numObjects + blockSize - 1) / blockSize); // TODO-1.2 - use the kernels you wrote to step the simulation forward in time. + kernUpdateVelocityBruteForce<<>>(numObjects, dev_pos, dev_vel1, dev_vel2); + kernUpdatePos <<>>(numObjects, dt, dev_pos, dev_vel1); // TODO-1.2 ping-pong the velocity buffers + glm::vec3 *tmp = dev_vel1; + dev_vel1 = dev_vel2; + dev_vel2 = tmp; } void Boids::stepSimulationScatteredGrid(float dt) { From 6c6a0c2ed930250c633fff1d8df092db0f601c6a Mon Sep 17 00:00:00 2001 From: HAN YAN Date: Wed, 9 Sep 2020 15:07:24 -0400 Subject: [PATCH 02/17] clamp speed --- src/kernel.cu | 22 +++++++++------------- 1 file changed, 9 insertions(+), 13 deletions(-) diff --git a/src/kernel.cu b/src/kernel.cu index 5fc5e13..fb4e79a 100644 --- a/src/kernel.cu +++ b/src/kernel.cu @@ -234,6 +234,8 @@ __device__ glm::vec3 computeVelocityChange(int N, int iSelf, const glm::vec3 *po glm::vec3 perceived_center = glm::vec3(0.0f, 0.0f, 0.0f); glm::vec3 c = glm::vec3(0.0f, 0.0f, 0.0f); glm::vec3 perceived_velocity = glm::vec3(0.0f, 0.0f, 0.0f); + int n1 = 0; + int n3 = 0; for (int i = 0; i < N; i++) { if (i == iSelf) { @@ -241,27 +243,24 @@ __device__ glm::vec3 computeVelocityChange(int N, int iSelf, const glm::vec3 *po } float dist = glm::distance(pos[i], pos[iSelf]); // Rule 1: boids fly towards their local perceived center of mass, which excludes themselves - int n1 = 0; if (dist < rule1Distance) { perceived_center += pos[i]; n1++; } - perceived_center /= n1; - + // Rule 2: boids try to stay a distance d away from each other if (dist < rule2Distance) { c -= (pos[i] - pos[iSelf]); } // Rule 3: boids try to match the speed of surrounding boids - int n3 = 0; if (dist < rule3Distance) { perceived_velocity += vel[i]; n3++; } - perceived_velocity /= n3; } - + perceived_center /= n1; + perceived_velocity /= n3; glm::vec3 new_vel = vel[iSelf] + (perceived_center - pos[iSelf]) * rule1Scale + c * rule2Scale + perceived_velocity * rule3Scale; @@ -282,13 +281,10 @@ __global__ void kernUpdateVelocityBruteForce(int N, glm::vec3 *pos, glm::vec3 new_vel = computeVelocityChange(N, index, pos, vel1); // Clamp the speed - new_vel.x = new_vel.x > maxSpeed ? maxSpeed : new_vel.x; - new_vel.y = new_vel.y > maxSpeed ? maxSpeed : new_vel.y; - new_vel.z = new_vel.z > maxSpeed ? maxSpeed : new_vel.z; - - new_vel.x = new_vel.x < -maxSpeed ? -maxSpeed : new_vel.x; - new_vel.y = new_vel.y < -maxSpeed ? -maxSpeed : new_vel.y; - new_vel.z = new_vel.z < -maxSpeed ? -maxSpeed : new_vel.z; + float speed = glm::length(new_vel); + if (speed > maxSpeed) { + new_vel = glm::normalize(new_vel) * maxSpeed; + } // Record the new velocity into vel2. Question: why NOT vel1? vel2[index] = new_vel; From 2d2a36815d05abf6fb55aa680fd8233939c535eb Mon Sep 17 00:00:00 2001 From: HAN YAN Date: Wed, 9 Sep 2020 15:46:09 -0400 Subject: [PATCH 03/17] 2.1 checkpoint1 --- src/kernel.cu | 24 ++++++++++++++++++++++++ 1 file changed, 24 insertions(+) diff --git a/src/kernel.cu b/src/kernel.cu index fb4e79a..a439404 100644 --- a/src/kernel.cu +++ b/src/kernel.cu @@ -169,6 +169,12 @@ 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!"); + cudaDeviceSynchronize(); } @@ -332,6 +338,12 @@ __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 + for (int i = 0; i < N; i++) { + glm::vec3 offset = glm::floor((pos[i] - gridMin) * inverseCellWidth); + int gridIndex = gridIndex3Dto1D(offset.x, offset.y, offset.z, gridResolution); + indices[i] = i; + gridIndices[i] = gridIndex; + } } // LOOK-2.1 Consider how this could be useful for indicating that a cell @@ -349,6 +361,16 @@ __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 prev = -1; + for (int i = 0; i < N; i++) { + if (particleGridIndices[i] != prev) { + if (prev != -1) { + gridCellEndIndices[prev] = i - 1; + } + prev = particleGridIndices[i]; + gridCellStartIndices[prev] = i; + } + } } __global__ void kernUpdateVelNeighborSearchScattered( @@ -439,6 +461,8 @@ void Boids::endSimulation() { cudaFree(dev_pos); // TODO-2.1 TODO-2.3 - Free any additional buffers here. + cudaFree(dev_particleArrayIndices); + cudaFree(dev_particleGridIndices); } void Boids::unitTest() { From 2492e276d48fd14ff60e040ec3e53e3706850f2d Mon Sep 17 00:00:00 2001 From: HAN YAN Date: Wed, 9 Sep 2020 19:05:29 -0400 Subject: [PATCH 04/17] part 2.1 checkpoint 2 --- src/kernel.cu | 52 +++++++++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 52 insertions(+) diff --git a/src/kernel.cu b/src/kernel.cu index a439404..8f8f8d0 100644 --- a/src/kernel.cu +++ b/src/kernel.cu @@ -387,6 +387,58 @@ __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 = (blockIdx.x * blockDim.x) + threadIdx.x; + if (index >= N) { + return; + } + float r = imax(imax(rule1Distance, rule2Distance), rule3Distance); + + glm::vec3 pos_min = glm::clamp(pos[index] - glm::vec3(r, r, r), -scene_scale, scene_scale); + glm::vec3 pos_max = glm::clamp(pos[index] + glm::vec3(r, r, r), -scene_scale, scene_scale); + glm::vec3 cell_min = glm::floor((pos_min - gridMin) * inverseCellWidth); + glm::vec3 cell_max = glm::floor((pos_max - gridMin) * inverseCellWidth); + + glm::vec3 perceived_center = glm::vec3(0.0f, 0.0f, 0.0f); + glm::vec3 c = glm::vec3(0.0f, 0.0f, 0.0f); + glm::vec3 perceived_velocity = glm::vec3(0.0f, 0.0f, 0.0f); + int n1 = 0; + int n3 = 0; + + for (int i = cell_min.x; i <= cell_max.x; i++) { + for (int j = cell_min.y; j <= cell_max.y; j++) { + for (int k = cell_min.z; k <= cell_max.z; k++) { + int cellIdx = gridIndex3Dto1D(i, j, k, gridResolution); + int startIdx = gridCellStartIndices[cellIdx]; + int endIdx = gridCellEndIndices[cellIdx]; + for (int p = startIdx; p <= endIdx; p++) { + int idx = particleArrayIndices[p]; + float dist = glm::distance(pos[index], pos[idx]); + if (dist < rule1Distance) { + perceived_center += pos[idx]; + n1++; + } + if (dist < rule2Distance) { + c -= (pos[idx] - pos[index]); + } + if (dist < rule3Distance) { + perceived_velocity += vel1[idx]; + n3++; + } + } + } + } + } + perceived_center /= n1; + perceived_velocity /= n3; + glm::vec3 new_vel = vel1[index] + (perceived_center - pos[index]) * rule1Scale + + c * rule2Scale + perceived_velocity * rule3Scale; + // Clamp the speed + float speed = glm::length(new_vel); + if (speed > maxSpeed) { + new_vel = glm::normalize(new_vel) * maxSpeed; + } + + vel2[index] = new_vel; } __global__ void kernUpdateVelNeighborSearchCoherent( From c89c3e017effa52042329bbfeb9125475f7f2045 Mon Sep 17 00:00:00 2001 From: HAN YAN Date: Wed, 9 Sep 2020 19:10:42 -0400 Subject: [PATCH 05/17] part 2.1 checkpoint 3 --- src/kernel.cu | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/src/kernel.cu b/src/kernel.cu index 8f8f8d0..3425b3f 100644 --- a/src/kernel.cu +++ b/src/kernel.cu @@ -175,6 +175,12 @@ void Boids::initSimulation(int N) { 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!"); + cudaDeviceSynchronize(); } @@ -515,6 +521,8 @@ void Boids::endSimulation() { // TODO-2.1 TODO-2.3 - Free any additional buffers here. cudaFree(dev_particleArrayIndices); cudaFree(dev_particleGridIndices); + cudaFree(dev_gridCellStartIndices); + cudaFree(dev_gridCellEndIndices); } void Boids::unitTest() { From b7050205cf1991bf35f13262cf5d23f692db6933 Mon Sep 17 00:00:00 2001 From: HAN YAN Date: Wed, 9 Sep 2020 20:05:10 -0400 Subject: [PATCH 06/17] part 2.1 checkpoint 4 --- src/kernel.cu | 19 ++++++++++++++++++- 1 file changed, 18 insertions(+), 1 deletion(-) diff --git a/src/kernel.cu b/src/kernel.cu index 3425b3f..6254ac6 100644 --- a/src/kernel.cu +++ b/src/kernel.cu @@ -486,13 +486,30 @@ void Boids::stepSimulationScatteredGrid(float dt) { // In Parallel: // - label each particle with its array index as well as its grid index. // Use 2x width grids. + dim3 fullBlocksPerGrid((numObjects + blockSize - 1) / blockSize); + 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. + dev_thrust_particleArrayIndices = thrust::device_ptr(dev_particleArrayIndices); + dev_thrust_particleGridIndices = thrust::device_ptr(dev_particleGridIndices); + 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 + // cell's data pointers in the array of boid indices + kernResetIntBuffer <<>> (gridCellCount, dev_gridCellStartIndices, -1); + 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, + dev_gridCellStartIndices, dev_gridCellEndIndices, dev_particleArrayIndices, + dev_pos, dev_vel1, dev_vel2); // - Update positions + kernUpdatePos <<>> (numObjects, dt, dev_pos, dev_vel1); // - Ping-pong buffers as needed + glm::vec3 *tmp = dev_vel1; + dev_vel1 = dev_vel2; + dev_vel2 = tmp; } void Boids::stepSimulationCoherentGrid(float dt) { From 51e7d1fbc6377b72ff3efb78d90cd1374939cb64 Mon Sep 17 00:00:00 2001 From: HAN YAN Date: Wed, 9 Sep 2020 20:45:29 -0400 Subject: [PATCH 07/17] 2.3 checkpoint 1 --- src/kernel.cu | 75 +++++++++++++++++++++++++++++++++++++++++++++++++++ src/main.cpp | 2 +- 2 files changed, 76 insertions(+), 1 deletion(-) diff --git a/src/kernel.cu b/src/kernel.cu index 6254ac6..9753317 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_coh_pos; +glm::vec3 *dev_coh_vel; // LOOK-2.1 - Grid parameters based on simulation parameters. // These are automatically computed for you in Boids::initSimulation @@ -181,6 +183,13 @@ void Boids::initSimulation(int N) { cudaMalloc((void**)&dev_gridCellEndIndices, gridCellCount * sizeof(int)); checkCUDAErrorWithLine("cudaMalloc dev_gridCellEndIndices failed!"); + // 2.3 + cudaMalloc((void**)&dev_coh_pos, N * sizeof(glm::vec3)); + checkCUDAErrorWithLine("cudaMalloc dev_coh_pos failed!"); + + cudaMalloc((void**)&dev_coh_vel, N * sizeof(glm::vec3)); + checkCUDAErrorWithLine("cudaMalloc dev_coh_vel failed!"); + cudaDeviceSynchronize(); } @@ -464,6 +473,59 @@ __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 = (blockIdx.x * blockDim.x) + threadIdx.x; + if (index >= N) { + return; + } + float r = imax(imax(rule1Distance, rule2Distance), rule3Distance); + + glm::vec3 pos_min = glm::clamp(pos[index] - glm::vec3(r, r, r), -scene_scale, scene_scale); + glm::vec3 pos_max = glm::clamp(pos[index] + glm::vec3(r, r, r), -scene_scale, scene_scale); + glm::vec3 cell_min = glm::floor((pos_min - gridMin) * inverseCellWidth); + glm::vec3 cell_max = glm::floor((pos_max - gridMin) * inverseCellWidth); + + glm::vec3 perceived_center = glm::vec3(0.0f, 0.0f, 0.0f); + glm::vec3 c = glm::vec3(0.0f, 0.0f, 0.0f); + glm::vec3 perceived_velocity = glm::vec3(0.0f, 0.0f, 0.0f); + int n1 = 0; + int n3 = 0; + + for (int k = cell_min.z; k <= cell_max.z; k++) { + for (int j = cell_min.y; j <= cell_max.y; j++) { + for (int i = cell_min.x; i <= cell_max.x; i++) { + int cellIdx = gridIndex3Dto1D(i, j, k, gridResolution); + int startIdx = gridCellStartIndices[cellIdx]; + int endIdx = gridCellEndIndices[cellIdx]; + for (int idx = startIdx; idx <= endIdx; idx++) { + float dist = glm::distance(pos[index], pos[idx]); + if (dist < rule1Distance) { + perceived_center += pos[idx]; + n1++; + } + if (dist < rule2Distance) { + c -= (pos[idx] - pos[index]); + } + if (dist < rule3Distance) { + perceived_velocity += vel1[idx]; + n3++; + } + } + } + } + } + + perceived_center /= n1; + perceived_velocity /= n3; + glm::vec3 new_vel = vel1[index] + (perceived_center - pos[index]) * rule1Scale + + c * rule2Scale + perceived_velocity * rule3Scale; + // Clamp the speed + float speed = glm::length(new_vel); + if (speed > maxSpeed) { + new_vel = glm::normalize(new_vel) * maxSpeed; + } + + vel2[index] = new_vel; } /** @@ -518,10 +580,19 @@ void Boids::stepSimulationCoherentGrid(float dt) { // In Parallel: // - Label each particle with its array index as well as its grid index. // Use 2x width grids + dim3 fullBlocksPerGrid((numObjects + blockSize - 1) / blockSize); + 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. + dev_thrust_particleArrayIndices = thrust::device_ptr(dev_particleArrayIndices); + dev_thrust_particleGridIndices = thrust::device_ptr(dev_particleGridIndices); + 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); + 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 @@ -540,6 +611,10 @@ void Boids::endSimulation() { cudaFree(dev_particleGridIndices); cudaFree(dev_gridCellStartIndices); cudaFree(dev_gridCellEndIndices); + + // 2.3 + cudaFree(dev_coh_pos); + cudaFree(dev_coh_vel); } void Boids::unitTest() { diff --git a/src/main.cpp b/src/main.cpp index b82c8c6..063a977 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -14,7 +14,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 From 477801cbc3a4872947c39d01d0444369178636b5 Mon Sep 17 00:00:00 2001 From: HAN YAN Date: Wed, 9 Sep 2020 21:12:20 -0400 Subject: [PATCH 08/17] 2.3 checkpoint 2 --- src/kernel.cu | 26 ++++++++++++++++++++++++++ src/main.cpp | 2 +- 2 files changed, 27 insertions(+), 1 deletion(-) diff --git a/src/kernel.cu b/src/kernel.cu index 9753317..109c2c1 100644 --- a/src/kernel.cu +++ b/src/kernel.cu @@ -388,6 +388,16 @@ __global__ void kernIdentifyCellStartEnd(int N, int *particleGridIndices, } } +__global__ void kernReorderData(int N, int *particleArrayIndices, + glm::vec3 *pos, glm::vec3 *vel, glm::vec3 *coh_pos, glm::vec3 *coh_vel) { + int index = (blockIdx.x * blockDim.x) + threadIdx.x; + if (index < N) { + int d = particleArrayIndices[index]; + coh_pos[index] = pos[d]; + coh_vel[index] = vel[d]; + } +} + __global__ void kernUpdateVelNeighborSearchScattered( int N, int gridResolution, glm::vec3 gridMin, float inverseCellWidth, float cellWidth, @@ -596,9 +606,25 @@ void Boids::stepSimulationCoherentGrid(float dt) { // - BIG DIFFERENCE: use the rearranged array index buffer to reshuffle all // the particle data in the simulation array. // CONSIDER WHAT ADDITIONAL BUFFERS YOU NEED + kernReorderData << > > (numObjects, dev_particleArrayIndices, dev_pos, dev_vel1, dev_coh_pos, dev_coh_vel); + glm::vec3 *tmp = dev_pos; + dev_pos = dev_coh_pos; + dev_coh_pos = tmp; + + tmp = dev_vel1; + dev_vel1 = dev_coh_vel; + dev_coh_vel = tmp; + // - Perform velocity updates using neighbor search + kernUpdateVelNeighborSearchCoherent << > > (numObjects, gridSideCount, gridMinimum, gridInverseCellWidth, gridCellWidth, + dev_gridCellStartIndices, dev_gridCellEndIndices, + dev_pos, dev_vel1, dev_vel2); // - Update positions + kernUpdatePos << > > (numObjects, dt, dev_pos, dev_vel1); // - Ping-pong buffers as needed. THIS MAY BE DIFFERENT FROM BEFORE. + tmp = dev_vel1; + dev_vel1 = dev_vel2; + dev_vel2 = tmp; } void Boids::endSimulation() { diff --git a/src/main.cpp b/src/main.cpp index 063a977..ddd0e3b 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -15,7 +15,7 @@ // LOOK-2.1 LOOK-2.3 - toggles for UNIFORM_GRID and COHERENT_GRID #define VISUALIZE 1 #define UNIFORM_GRID 1 -#define COHERENT_GRID 0 +#define COHERENT_GRID 1 // LOOK-1.2 - change this to adjust particle count in the simulation const int N_FOR_VIS = 5000; From 67ed894f24bd1a6e0b9d743e90719715ef02bf02 Mon Sep 17 00:00:00 2001 From: Tracy-Yan <33076509+Tracy-Yan@users.noreply.github.com> Date: Sat, 12 Sep 2020 16:43:02 -0400 Subject: [PATCH 09/17] Add files via upload --- images/cap1.gif | Bin 0 -> 1733598 bytes 1 file changed, 0 insertions(+), 0 deletions(-) create mode 100644 images/cap1.gif diff --git a/images/cap1.gif b/images/cap1.gif new file mode 100644 index 0000000000000000000000000000000000000000..cd4f69a305e4f93d486b03b640bcf59be2e9081d GIT binary patch literal 1733598 zcmWifXHe4*566F9&@#*3SoVg3h($olD0>uSl%=4kh%5mWmC=?01qzhCv20YpGQ^^Y zW$(Qa6crT^6*u7E`9H6c%jLe8yW~akxr>#JCC1QS81@l-0|4X>xgBw?r_D|AI6ZAN z1O)sS+J^$u|2dlf`N03vlK->3+)Dh4eEcdZihO)3yei5nNHtA8B)_$wU?iWGJHJl2 z8n-=~KTHkds*Mg&@>S%+JMr3v3rE-sT)T19*72xA@X=srbV#@rKc53!-(6SGPG9G` zDF0QYUbv1>C>rg;?PMY!)mtP>W@XB*CcHA34VLOaC4h0oqj&+ZDpeVAT|BTtx< zZqzlw;4p1#8ykD)5d0M9s-#?oF)uq;d1PtHB{SV6^%0&)cV)Ij zWYR*i8^as3!ke2NGA0}|Hys+rPB(2`$r%mLdlz

dJ#{@A=ZR^NTcL-VAB>S`Jj5t4%!44+KIQ^_-ZR;h7HnSx(xSR|PXdbsjkK zGf!(VuL?GW>0~-NX4!|NML4EKL^L_#nWsaUVNNX(VN23S=4H+-_@EZgdX>tQmdY%r zJJK_pA7n&iWP~#xIHk9`KV)9XY`Jp3+;?t4?#{xI<$2$V1;3@@k|P-z&a}+%%qH9P zrm&WDrw8=VmIrv|gAm5NLg~EUT&3(7-Dxs2bb{$P!3^D^pWaFf*=%&#Vumj)qL$~+ zXJn+NWj3WX(bJk+?xtt8rnA}}&>9*aWYEUxshcS+bLp%%P3aR2S)0s9Bh9rRs%9Bk zW3xF!Z)n>K^zCit$V|iJyTAKM$2N|q+a z()TXu8fx;KDy*eC;>Mp@1aLL%V)lXMh=cxORwB4Pp!&lj+kF;(( z-T}BA*G3-RUwq)#n<``1-mu&nI2+V9vhW^gM0~L+v3nHi`!M?3_S)#92QQxyA>6X| zIgM^45Ym;HhL*y_&=st@>{u;tM+=W$c`??}`hFzo{s9kr&+l>){nY(4|M*<>eCdE#@43zLlTL>TrH#d3f@s|rM%`ELD09y9aJ&Kj9x)N&e7$2HZSwnP z!eR;K@OJdy?!AN|j5J!O?AG`1A0`!`Zr+{G9HtM~9=uy^zfS2k{<3EJk0F27>hJtJ zPV-+&^gE`%bvH^OYlP&on-iwG>!*fjyMDKYl`Z%YoRZL&AW1J_e+xuX;1u=!bd{{E z&II8NNPlGLMX>SD%{Q4?j>`hTj5SA*zGs=@IEeNtlBnPHw^qo#$8^A6hOhu~JU@>v zpr2AjxB8;PjcZ>jPp~3;5`NWmD3fx0oaN0w6|#BF*nxLNiZ1@kxpaqf?G3s7BjpXh z!v22iy9%$z5VO?7)ziK38}iV1`CN_GS(RSJ=JG*jyE^EV6FZUoTodc(Gma&nz~7Ba zgdQ^2O=I#oMqu)}R2t+4HXl(fq;^nYUi!K!E&6b^_BQxg4Trg{|Jr6|Ni;= z^NrLAuB=n(o4=jZf-5W0z1$aE29uJNQv>^T;zsxM}7zt{eJ&unzG&|`4#(T{Vk z9w9&GZAka&3$}VT+LrI%+^Bl4Kw;)zkX^PL{8bS zLUPw98Cv?KFY8`1f}butwOdg;VTdHJc=~fgest(u_;fD+_n^|ApSq^Ft2;sm+U(z;cNN9?JYAlx|WPhsC@ z_52jr8ie~fvPkcYxM`})N7Nt^Nqmjw&xVC0ek8b-vIQ(p?Y)68-@Jw|e^(A*@|eB* z?Zh!}Woo+;a%A_(9p&T0!;W6>3>e?@ZGK)tx@1F-oC0|-4Ty@xEWPrLP)!(5>NVh) z$uQsPWpL?@sHBSNcl(P9?5mc~=k$Cv{f1D?$s^PvuX-@Bu@Dh)r z?#S>f;whz*cNTA-z7cSZU!qee@)hl(6aos?)FV1%`P@;8+|RVTIFp0oy=!fCAfh0# zrC6(9V?Cs#?N&aj!dIWI@d|0@ncbNGv>K(WZvE)sl0sA<3Z#mT)ern#i zo5wcu4z$Dqtp&3fu3$*18i}~8fQTzd05l!*Lg%V=*XAmum;;n_r1Y2F-33Tv+~Z5f zJ~^AyBt-5W(Qr2SJ#@rD|4pQ76+@eIL?&H=68X$$_U;U{@77#iLl*hI7b1$xY3p># zN@H{X`s1?k+D-BirMrPvb}KXglxlhq!;Al>T_SBR>+4-h3ND$Zc^=l*p7H+r?iBUF zJyd#xzPI%;46!k2=F{z}50L{ z^xhsD@!*mRG9P|m+2)O1+Sf0dkmQ!H+`I2iOX;>q66ZG?zPa$mxQGQdGYla9O0j=r z;p-fj=-F?t56W@jeIULiCmZ&dYYFH^3k1#6Erf5uNF z-@J#^lTe&Iy{2$wWv=GmB>C*vw)>NVd7Z5wRek`--P4h?i})B89)lb#oU6vYzG@Ka zZaK@YIWpi5Z_)ULWnl*jRPo}cLemD$9iCsg+u#1(oyx!EFQxqtKa`aJbSGZ(cH&?4 zUk~qo>`ABGr;4^M;+rNgMjF@=9NR1YTyFQ1HdrbFH37C{=B z1X=e2`zLC8+4}=GCcc&b_EA?jVm+i25m}&699)p+Fzq(`@s82oEqvyIJ-K|O3hJ{> zn0&nOXyQjR-o(b2|H;yz4ZKSf^}UQ90k z)%E-DXX3NpFCb2PJyQR6&mA#-EjGC~VAM6A8sE90;_r2cf>9)t@!NQY&KbV=~o zcJA%okW<&!g2_Mcj1KpwGP}OlIsM)CK0H{O?E3j=^6#hHhriePpB?TEI{o`va`Dj^Jv_#VqHkFWl`*-6HJ#x!EzwoDT*s0 z(LFkmP?+e&O7x+MI7N$tK`{NP!~jCl#ptBK!lWQp((bRsN|JcnKvD!IIf{^cJvup> zKnaCQ-&{(L+e;>jQ^^?W??=fVrc`PnHI+rBEm0E=1F5H9Q?oHCxrCJb=#&%ilyeCw zrAsO0dnpX@)WfOEHKh*fICUMf>Otp$@t+)QnYPGc3OJ!Yj@1*UcG zrLo0nJs6tUW*UG`?JJ~>uxMjTwD*^2Lz}c2O!^!lePPLMKm|D+oxZx1{$elvK0JLL zld(a_c(+M?9i6ew%Gg=TxEx3Wt!I42WPT@P-YQK0S(tgi%KQ^Z9$3l*B%-j z2oHi5(YZS4JVkK>7=TwIOVBV&7!i9&$PzEgk{qUsb!5rzXQ2!;WhAl{J+qZ>rYqdc zR`1Bx+)q2i_mW!lxvqDsSJ0h^-go5l5NnYBg54a>!ESj9~eB`t>O zOv94)n`h^r%l@9*f4cuxpYdREPk(a#r3yM4UTuG?!gaYCZ(AiPS>u{orRY^d7^(6ouJN(GgIKBY zD=s&cs0~oB_4cX_94Rv_t_@BtzOqsqzFib1Sr_G16zNqLJyLk1xGv^I;q8?=;z&Wf zWIg3XLA5Gx9u@$QnTdxh^$j+uoQ!(Tt@_+s#W=6Kg)2qXEN%_}a4EdY2r8;3aF#v> z_$BYvrxsT4K}%2Ey?byk*(P0Fvq>#Wz27NWn6Pr6otpDZvY}5hr`M}t z=w9|N&FKmL36Jt8jJgM8whJnq>qom`=mcipZBh33b zn-4Chw#R3VwO%T2nLQ4*Ahzy_K`_|(Y$C$)d4j%F zBZ$$4qXG!zLv75%wR;d#zlUIdL^Zx;t`4ldE~G{W4~l{1cjL0xTPF(u&-)KGm#U(k z$Jioae%yd53VPd)>-?2i?eKO;7Fc6l=)avuajN~33ZiH#hMNiW4go?MmCtWI8Ww}y zN{ewTiGgC<*E~Sy3t3AwUl29G z-fHUFj z@w(Q1m7Wh)&yI_C8=dH`x`}{v#vk|No~r8s69AFY7|n(_m2mLbuMU-9NC3KrOPu2& zGS&qRlZxqDIbQqQlj~?heESsl@g45Zqj9z2Tn?S_Dh$LNCs)n7kbxh>Q<~_z3%jHl zTgUD2}id*c(WA zyCi(CG`75ycORL22H+Sg8FM%n=*wcOXSw7 zP&bKdUX6KXPn@QS+#Ka*?{=heL9P5AdZT%hH35Z>v39GEKYbW}S=WIX<=a8d0jkrB z1w6ydfhRoqC!WrFEX^MG1gFx;E_AL9tBLJPT+PI3k@}7iIuH6b)EYVa?l>&eI2Pm( zJF5w}dBZ%3;$9&{(E9w0SXJj3P~=H%$qM7l7SC+a5qi(oSv$TSBM^Jz>Gz!R>voTF7f zWM3oyXLh&_ysjB{Jzc~vB=$=9%1cfzHe+Rs6ZZVjWU|3N-ruiRQFDsmya4H*0DtLX zGMIhkDczhvSJ!z@o5cCPh;hO`ANAns#{xe1w*iN|lkyn zyJHDE;55MTeWb^G#HYC{6buJ~=r8&q9&v)?ZXQ!ao;0`Z6?Jf#pM$zI*5m zPh`i=uD(}d!U_JcaBLjMNn~aChpD=El7TNdqc^!PZw-x8I3$Ywy!vO?;&x@@!`5Q1 zX2fhT`8djAP9jGjy6}W`p+q`-<@k7KY)qQpr;{MgYuGr043Wl`z&7WWf+%o8ln-@^%Ro!c>wu`R1sD_b(#-;%k)$72q=hSeHlPBV7F zy6|Ka5kvpYq9HsMLZAHo zCr+Zgzsm2&o5e2xu;bkoqeuih3~|PV|GYL{a^TGN*KX7jKv%L&J_Jac`#9nGl!#daWA{5VvNjnlBUw-hG6Oyu$XcZ2HIFF7T zDXF+qtKaoz#0tWE5PSC!{d>`aYXCp!w!VF`dl937fIW$I!vBF@{&#WYtC-Yx5N35G ze$!UtaLNki8cu%Fxv|5EbXo#XgsF8?3amiWW@+k`IaNRj3`hFC#?d5>g>EcOzpQ1wdt92D+8W=%y} zs(U283&z>0njQ~@+9wWOTpv_A+mWH=R8ty*ZAsG1%-ZI6le6rpa;z2Z<*%G8bqcc z{Mo3XB4Yi7UZymHUG4C5aclm;rBes`+j%e%eCXMi|6&vT!(U|%{kbT9w%krfC;88t zF!f`BcJ*mGP<_?=jq9wc3b{%u*d)QPJ`uH>tFFywjOH2c0ah@dQFhjb;W$z6sOgGd z>ybda*p)VCZSUx-yOt%;aP3;_i)z|U)|P4$IW7SG9G8I1!7EN#-jRk19m%LbTj0}f z6$7KJ2%oLP8o8|7jxp*KP$3vf7|_|T_jxwkVH>2%T;5RYjxMfR?d%0S>*r(r`*~kO{lZrz$mRvKX9A|Mg#!*7j6kv@O=>GC_4Q+ z>c>2F*MisG>&nACO{|sVw+bb@bcC?~1X`G=JAQW%I0w^08TnFef+}U@KXh~Ii(Gw^ zE=Ok8Lvh=0$Z5ftnQ$WIOthRr#4Fpp3E!}5i|(Flw{D!d_2B}q%UN9Z`HTd2+)swD zLdpd^ne3eGW}~1$O+QX#wV`k?q7&e9nl61=2jC&)JD_)P#YIZb5_6UJ5)x7>t~r zl9&s6)bM2|@8-pv*CA;aOCuMSYc=(ihH_7f4|ynN&MIxbrC6rmAViZ70R?G_96ia0 z2^^?vo9S^!)e=Ruv%#V~FqDUD6;801Cs%6PJTu{<_|JI-d9BOh=a7vv(T-HIsa0rA{ z#UWk&Dp0A-yaGz{C&j=FQBV(D@|Xfd@?3z!Vl%wnCmceJ0igWOAo-Uy`bH<9&W_cJ zUip|R+lzvykVJW)8V)+TN`$#7+KI9ki?uAP&Fa|r-4=-O7X1PJnfGjNDCi;GPcDt3 zN>M;tBKb8hs%Vu`b#kxvQNyr*3q3U4Zhdj^eJLXfn-2Z{Jt0j?e zgtcF(dA+`t6hYBUJsgy<(<+DD1e=L66R1=g%4o@y*9QPZ19DO6)7D^h+ks*{S@2g8 z?nG+qi9(K_irON<{&jjMiY0Z+k9qd$M30OdmM|(sWTx*~xZLEuu6j+;FgVCO?iB&3 zmgDUNS6Nu22Fd8)G8vHBa+sE#rNN*z7A(b)o@MF04~A8ZP| zxuydkO;f?u{&k4n5h)b=MSA!zZLvTJv&|AD+?3M5zXe7Cpq77N)W8veaEneLX~o)1 zK#66OC?rThSrCTwQ$Eqt|5S3eAHEk1y_Y6DaHU&QSdlkG9((?+T8&1@hR6@iG>G~X zT{z7V*d4d7_0eSTMKMsi$`dpYvx&S%&P3d8GG`OdC*%?b4Fs;5lJU(7oW+XMux1AD zc(K7MUi*&twx#v-=Bpn`JW`I$HG?ecLxxL5>G8fdPZ}y8l)U-OKqKydYho znElpTjYZ2UxB#9B@`TiE`Ot0N3&2hqS__pjncIc7b5$9Q7WgSRWeJ$)^(QGPAW0OO zv%DIi+cZEcBpIb8aX5^v)%)E3iA==a=FI}}G5&@>K}x0dp-#cf^F2c zRA111V=L4gVcoM+Z2A2imk*W-RwU9dnOVG)dL;??jc1WmmQcSwSj#D!US-7B4w