diff --git a/.gitignore b/.gitignore index b699c9e..019f727 100644 --- a/.gitignore +++ b/.gitignore @@ -562,3 +562,16 @@ xcuserdata *.xccheckout *.moved-aside *.xcuserstate +/build/ +/bin/ +/lib/ +/x64/ +/Debug/ +/Release/ +/CMakeFiles/ +/CMakeCache.txt +/cmake_install.cmake +/.vs/ +/*.user +/*.vcxproj.user +/*.pdb diff --git a/README.md b/README.md index ee39093..639bbe5 100644 --- a/README.md +++ b/README.md @@ -1,11 +1,120 @@ +**Note:** Using 1 late day + extension granted by Mr. Mohammed + **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) -### (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.) +* **Cindy Wei** +* **Tested on:** + - **OS:** Windows 11 Home (Version 23H2, Build 22631.3447) + - **CPU:** Intel Core i7-13700HX @ 2.10 GHz (20 cores, 28 threads) + - **RAM:** 16 GB DDR5 + - **GPU:** NVIDIA GeForce RTX 4070 Laptop GPU (8 GB VRAM) + - **Machine:** Lenovo Legion Pro 5 16IRX9 (Personal) + +## Demo + +![Naive 5k Boids](images/part1.gif) +*Naive implementation with 5,000 boids* + +![Scattered 5k Boids](images/part2.gif) +*Scattered implementation with 5,000 boids* + +![Coherent 5k Boids](images/part3.gif) +*Coherent implementation with 5,000 boids* + +## Summary + +I implemented a 3D flocking simulation (Reynolds Boids) on the GPU with three execution modes: +- **Naive**: All-pairs neighbor search +- **Uniform Grid (scattered)**: Grid indices sorted; boid arrays left in original order +- **Uniform Grid (coherent)**: Grid indices sorted; boid positions/velocities reordered into cell order for more coherent memory access + +I analyzed performance vs. boid count, the impact of coherent layout, block size, and visualization overhead. + +## Toggling Modes + +In `src/main.cpp`, use the provided defines/toggles: +- `NAIVE_MODE` +- `SCATTERED_MODE` +- `COHERENT_MODE` +- `VISUALIZATION` (0 = off/headless timing, 1 = on) + +## Implementation Details + +### Part 1 — Naive Boids +Each timestep, every boid checks all others and applies: +- **Cohesion** (center of mass) +- **Separation** (short-range repulsion) +- **Alignment** (match neighbors' velocity) + +**CUDA kernels:** +- Velocity update (neighbor rules) +- Position integration (with simple bounds handling) + +**Code locations:** +- `src/main.cpp` - app wiring, toggles, timing, GL +- `src/kernel.cu` - device structs, kernels, host launchers +- Search for `TODO-1.2` / `LOOK-1.2` for implementation details + +### Part 2 — Uniform Grid Acceleration + +#### 2.1 Scattered Layout +- Compute each boid's grid cell index +- `Thrust::sort_by_key(particleGridIndex, particleArrayIndex)` +- Parallel sweep → `gridCellStart/End` arrays +- Neighbor kernel iterates adjacent cells only +- Position/velocity arrays stay in original order + +#### 2.3 Coherent Layout +- After sorting, reorder positions/velocities into `posCoherent/velCoherent` +- Neighbor kernel reads contiguous ranges directly +- Improves memory coalescing and cache locality + +**Grid configuration:** `cellw = neighbor radius` → 27 cells in 3D (`nb=27`) + +## Performance Results + +### Framerate vs Number of Boids +![FPS vs N](images/framerate_vs_num_boids.png) + + +### Block Size Optimization +![Optimal performance at 128-512 threads/block](images/block_size.png) +*Performance improvement with coherent memory access* + + + +### Performance Summary (nb=27, cellw=10.00, block=128) + +| N (boids) | vis | Scattered FPS | Coherent FPS | Improvement | +|-----------|-----|---------------|--------------|-------------| +| 2,000 | 0 | 8,500 | 9,650 | +13.5% | +| 5,000 | 0 | 5,000 | 5,400 | +8.0% | +| 10,000 | 0 | 4,350 | 4,850 | +11.5% | + +## Performance Analysis + +### Q1: How does boid count affect performance? +- **Naive**: O(N²) complexity, memory-bound at scale +- **Scattered grid**: Near O(N) with sorting overhead +- **Coherent grid**: Same complexity but better constants due to memory coalescing + +### Q2: How do block size/count affect performance? +- **Optimal**: 128-256 threads/block +- **Too small**: Low occupancy, poor latency hiding +- **Too large**: Register pressure limits occupancy + +### Q3: Does coherent layout improve performance? +**Yes**, +5-15% improvement expected due to: +- Contiguous memory access patterns +- Better cache utilization +- Reduced memory divergence + +### Q4: Cell width impact (27 vs 8 cells)? +- **27 cells**: More cell iterations, fewer boids per cell +- **8 cells**: Fewer iterations, more boids per cell +- **Result**: 27 cells generally faster for N ≥ 5k due to better spatial locality + + diff --git a/images/block_size.png b/images/block_size.png new file mode 100644 index 0000000..5fc783d Binary files /dev/null and b/images/block_size.png differ diff --git a/images/blocksize_output.png b/images/blocksize_output.png new file mode 100644 index 0000000..f7852ed Binary files /dev/null and b/images/blocksize_output.png differ diff --git a/images/framerate_vs_num_boids.png b/images/framerate_vs_num_boids.png new file mode 100644 index 0000000..7df2c09 Binary files /dev/null and b/images/framerate_vs_num_boids.png differ diff --git a/images/part1.gif b/images/part1.gif new file mode 100644 index 0000000..72adb22 Binary files /dev/null and b/images/part1.gif differ diff --git a/images/part2.gif b/images/part2.gif new file mode 100644 index 0000000..c99bd29 Binary files /dev/null and b/images/part2.gif differ diff --git a/images/part3.gif b/images/part3.gif new file mode 100644 index 0000000..1177511 Binary files /dev/null and b/images/part3.gif differ diff --git a/src/kernel.cu b/src/kernel.cu index 7149917..0c745f0 100644 --- a/src/kernel.cu +++ b/src/kernel.cu @@ -8,6 +8,11 @@ #include #include #include +#include +#include +#include +#include + #include #include @@ -25,6 +30,15 @@ #define imin( a, b ) ( ((a) < (b)) ? (a) : (b) ) #endif +#ifndef NEIGHBOR_MODE +#define NEIGHBOR_MODE 27 +#endif + +#ifndef CELL_WIDTH_MULT +#define CELL_WIDTH_MULT 2.0f +#endif + + #define checkCUDAErrorWithLine(msg) checkCUDAError(msg, __LINE__) /** @@ -47,7 +61,7 @@ void checkCUDAError(const char *msg, int line = -1) { *****************/ /*! Block size used for CUDA kernel launch. */ -#define blockSize 128 +#define blockSize 512 // LOOK-1.2 Parameters for the boids algorithm. // These worked well in our reference implementation. @@ -71,7 +85,12 @@ void checkCUDAError(const char *msg, int line = -1) { int numObjects; dim3 threadsPerBlock(blockSize); -// LOOK-1.2 - These buffers are here to hold all your boid information. +int Boids::getThreadsPerBlock() { + return static_cast(threadsPerBlock.x); +} + + +// LOOK-1.2 - These buffers are here hold all your boid information. // These get allocated for you in Boids::initSimulation. // Consider why you would need two velocity buffers in a simulation where each // boid cares about its neighbors' velocities. @@ -95,7 +114,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_posCoherent = nullptr; +glm::vec3* dev_vel1Coherent = nullptr; +glm::vec3* dev_vel2Coherent = nullptr; // LOOK-2.1 - Grid parameters based on simulation parameters. // These are automatically computed for you in Boids::initSimulation int gridCellCount; @@ -143,6 +164,20 @@ __global__ void kernGenerateRandomPosArray(int time, int N, glm::vec3 * arr, flo } } +__global__ void kernReorderDataCoherent( + int N, + const int* particleArrayIndices, + const glm::vec3* posIn, + const glm::vec3* vel1In, + glm::vec3* posOut, + glm::vec3* vel1Out) { + int s = blockIdx.x * blockDim.x + threadIdx.x; + if (s >= N) return; + int i = particleArrayIndices[s]; + posOut[s] = posIn[i]; + vel1Out[s] = vel1In[i]; +} + /** * Initialize memory, update some globals */ @@ -167,7 +202,7 @@ void Boids::initSimulation(int N) { checkCUDAErrorWithLine("kernGenerateRandomPosArray failed!"); // LOOK-2.1 computing grid params - gridCellWidth = 2.0f * std::max(std::max(rule1Distance, rule2Distance), rule3Distance); + gridCellWidth = CELL_WIDTH_MULT * std::max(std::max(rule1Distance, rule2Distance), rule3Distance); int halfSideCount = (int)(scene_scale / gridCellWidth) + 1; gridSideCount = 2 * halfSideCount; @@ -179,6 +214,22 @@ void Boids::initSimulation(int N) { gridMinimum.z -= halfGridWidth; // TODO-2.1 TODO-2.3 - Allocate additional buffers here. + + cudaMalloc(&dev_particleArrayIndices, numObjects * sizeof(int)); + cudaMalloc(&dev_particleGridIndices, numObjects * sizeof(int)); + dev_thrust_particleArrayIndices = thrust::device_pointer_cast(dev_particleArrayIndices); + dev_thrust_particleGridIndices = thrust::device_pointer_cast(dev_particleGridIndices); + + cudaMalloc(&dev_gridCellStartIndices, gridCellCount * sizeof(int)); + cudaMalloc(&dev_gridCellEndIndices, gridCellCount * sizeof(int)); + checkCUDAErrorWithLine("grid allocs failed"); + + cudaMalloc(&dev_posCoherent, numObjects * sizeof(glm::vec3)); + cudaMalloc(&dev_vel1Coherent, numObjects * sizeof(glm::vec3)); + cudaMalloc(&dev_vel2Coherent, numObjects * sizeof(glm::vec3)); + + + cudaDeviceSynchronize(); } @@ -243,9 +294,52 @@ __device__ glm::vec3 computeVelocityChange(int N, int iSelf, const glm::vec3 *po // 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); + const glm::vec3 myPos = pos[iSelf]; + + glm::vec3 perceived_center(0.0f); + glm::vec3 c_sep(0.0f); + glm::vec3 perceived_velocity(0.0f); + + int cnt1 = 0; + int cnt3 = 0; + + for (int j = 0; j < N; ++j) { + if (j == iSelf) continue; + + glm::vec3 otherPos = pos[j]; + float dist = glm::length(otherPos - myPos); + + if (dist < rule1Distance) { + perceived_center += otherPos; + cnt1++; + } + if (dist < rule2Distance) { + c_sep -= (otherPos - myPos); + } + if (dist < rule3Distance) { + perceived_velocity += vel[j]; + cnt3++; + } + } + + glm::vec3 dv(0.0f); + + if (cnt1 > 0) { + perceived_center /= (float)cnt1; + dv += (perceived_center - myPos) * rule1Scale; + } + + dv += c_sep * rule2Scale; + + if (cnt3 > 0) { + perceived_velocity /= (float)cnt3; + dv += perceived_velocity * rule3Scale; + } + + return dv; } + /** * TODO-1.2 implement basic flocking * For each of the `N` bodies, update its position based on its current velocity. @@ -255,6 +349,17 @@ __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 i = threadIdx.x + blockDim.x * blockIdx.x; + if (i >= N) return; + + glm::vec3 dv = computeVelocityChange(N, i, pos, vel1); + glm::vec3 v = vel1[i] + dv; + + float speed = glm::length(v); + if (speed > maxSpeed) { + v = (v / speed) * maxSpeed; + } + vel2[i] = v; } /** @@ -270,7 +375,6 @@ __global__ void kernUpdatePos(int N, float dt, glm::vec3 *pos, glm::vec3 *vel) { glm::vec3 thisPos = pos[index]; thisPos += vel[index] * dt; - // Wrap the boids around so we don't lose them thisPos.x = thisPos.x < -scene_scale ? scene_scale : thisPos.x; thisPos.y = thisPos.y < -scene_scale ? scene_scale : thisPos.y; thisPos.z = thisPos.z < -scene_scale ? scene_scale : thisPos.z; @@ -282,6 +386,9 @@ __global__ void kernUpdatePos(int N, float dt, glm::vec3 *pos, glm::vec3 *vel) { pos[index] = thisPos; } + + + // LOOK-2.1 Consider this method of computing a 1D index from a 3D grid index. // LOOK-2.3 Looking at this method, what would be the most memory efficient // order for iterating over neighboring grid cells? @@ -292,6 +399,11 @@ __device__ int gridIndex3Dto1D(int x, int y, int z, int gridResolution) { return x + y * gridResolution + z * gridResolution * gridResolution; } +__device__ __forceinline__ int clampi(int v, int lo, int hi) { + return (v < lo) ? lo : ((v > hi) ? hi : v); +} + + __global__ void kernComputeIndices(int N, int gridResolution, glm::vec3 gridMin, float inverseCellWidth, glm::vec3 *pos, int *indices, int *gridIndices) { @@ -299,12 +411,24 @@ __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 i = blockIdx.x * blockDim.x + threadIdx.x; + if (i >= N) return; + + glm::vec3 p = pos[i]; + int gx = clampi(int((p.x - gridMin.x) * inverseCellWidth), 0, gridResolution - 1); + int gy = clampi(int((p.y - gridMin.y) * inverseCellWidth), 0, gridResolution - 1); + int gz = clampi(int((p.z - gridMin.z) * inverseCellWidth), 0, gridResolution - 1); + + gridIndices[i] = gridIndex3Dto1D(gx, gy, gz, gridResolution); + indices[i] = i; + + } // LOOK-2.1 Consider how this could be useful for indicating that a cell // does not enclose any boids __global__ void kernResetIntBuffer(int N, int *intBuffer, int value) { - int index = (blockIdx.x * blockDim.x) + threadIdx.x; + int index = blockIdx.x * blockDim.x + threadIdx.x; if (index < N) { intBuffer[index] = value; } @@ -316,24 +440,100 @@ __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 i = blockIdx.x * blockDim.x + threadIdx.x; + if (i >= N) return; + + int gi = particleGridIndices[i]; + + if (i == 0 || gi != particleGridIndices[i - 1]) { + gridCellStartIndices[gi] = i; + } + if (i == N - 1 || gi != particleGridIndices[i + 1]) { + gridCellEndIndices[gi] = i; + } } + + + + + + + __global__ void kernUpdateVelNeighborSearchScattered( - int N, int gridResolution, glm::vec3 gridMin, - float inverseCellWidth, float cellWidth, - 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. - // - 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. - // - 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 N, int gridResolution, glm::vec3 gridMin, + float inverseCellWidth, float cellWidth, + 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. + // - 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. + // - 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 sIdx = blockIdx.x * blockDim.x + threadIdx.x; // index into *sorted* arrays + if (sIdx >= N) return; + + int i = particleArrayIndices[sIdx]; + glm::vec3 p = pos[i]; + + int gx = clampi(int((p.x - gridMin.x) * inverseCellWidth), 0, gridResolution - 1); + int gy = clampi(int((p.y - gridMin.y) * inverseCellWidth), 0, gridResolution - 1); + int gz = clampi(int((p.z - gridMin.z) * inverseCellWidth), 0, gridResolution - 1); + + const int r = 1; + + glm::vec3 perceived_center(0.0f), c(0.0f), perceived_velocity(0.0f); + int cnt1 = 0, cnt3 = 0; + + for (int dz = -r; dz <= r; ++dz) { + int z = gz + dz; if (z < 0 || z >= gridResolution) continue; + for (int dy = -r; dy <= r; ++dy) { + int y = gy + dy; if (y < 0 || y >= gridResolution) continue; + for (int dx = -r; dx <= r; ++dx) { + int x = gx + dx; if (x < 0 || x >= gridResolution) continue; + + int cell = gridIndex3Dto1D(x, y, z, gridResolution); + int start = gridCellStartIndices[cell]; + if (start == -1) continue; + int end = gridCellEndIndices[cell]; + + for (int k = start; k <= end; ++k) { + int j = particleArrayIndices[k]; + if (j == i) continue; + + glm::vec3 pj = pos[j]; + float dist = glm::distance(p, pj); + + if (dist < rule1Distance) { perceived_center += pj; ++cnt1; } + if (dist < rule2Distance) { c -= (pj - p); } + if (dist < rule3Distance) { perceived_velocity += vel1[j]; ++cnt3; } + } + } + } + } + + glm::vec3 dv(0.0f); + if (cnt1 > 0) { + perceived_center /= (float)cnt1; + dv += (perceived_center - p) * rule1Scale; + } + dv += c * rule2Scale; + if (cnt3 > 0) { + perceived_velocity /= (float)cnt3; + dv += perceived_velocity * rule3Scale; + } + + glm::vec3 v = vel1[i] + dv; + float speed = glm::length(v); + if (speed > maxSpeed) v = (v / speed) * maxSpeed; + vel2[i] = v;; } + __global__ void kernUpdateVelNeighborSearchCoherent( int N, int gridResolution, glm::vec3 gridMin, float inverseCellWidth, float cellWidth, @@ -351,6 +551,61 @@ __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 i = blockIdx.x * blockDim.x + threadIdx.x; + if (i >= N) return; + + glm::vec3 p = pos[i]; + + int gx = clampi(int((p.x - gridMin.x) * inverseCellWidth), 0, gridResolution - 1); + int gy = clampi(int((p.y - gridMin.y) * inverseCellWidth), 0, gridResolution - 1); + int gz = clampi(int((p.z - gridMin.z) * inverseCellWidth), 0, gridResolution - 1); + + const int r = 1; + + glm::vec3 perceived_center(0.0f), c(0.0f), perceived_velocity(0.0f); + int cnt1 = 0, cnt3 = 0; + + for (int dz = -r; dz <= r; ++dz) { + int z = gz + dz; if (z < 0 || z >= gridResolution) continue; + for (int dy = -r; dy <= r; ++dy) { + int y = gy + dy; if (y < 0 || y >= gridResolution) continue; + for (int dx = -r; dx <= r; ++dx) { + int x = gx + dx; if (x < 0 || x >= gridResolution) continue; + + int cell = gridIndex3Dto1D(x, y, z, gridResolution); + int start = gridCellStartIndices[cell]; + if (start == -1) continue; + int end = gridCellEndIndices[cell]; + + for (int k = start; k <= end; ++k) { + if (k == i) continue; + glm::vec3 pj = pos[k]; + float dist = glm::distance(p, pj); + + if (dist < rule1Distance) { perceived_center += pj; ++cnt1; } + if (dist < rule2Distance) { c -= (pj - p); } + if (dist < rule3Distance) { perceived_velocity += vel1[k]; ++cnt3; } + } + } + } + } + + glm::vec3 dv(0.0f); + if (cnt1 > 0) { + perceived_center /= (float)cnt1; + dv += (perceived_center - p) * rule1Scale; + } + dv += c * rule2Scale; + if (cnt3 > 0) { + perceived_velocity /= (float)cnt3; + dv += perceived_velocity * rule3Scale; + } + + glm::vec3 v = vel1[i] + dv; + float speed = glm::length(v); + if (speed > maxSpeed) v = (v / speed) * maxSpeed; + vel2[i] = v; + } /** @@ -359,6 +614,17 @@ __global__ void kernUpdateVelNeighborSearchCoherent( 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 + dim3 fullBlocksPerGrid((numObjects + blockSize - 1) / blockSize); + + kernUpdateVelocityBruteForce << > > ( + numObjects, dev_pos, dev_vel1, dev_vel2); + checkCUDAErrorWithLine("kernUpdateVelocityBruteForce failed!"); + + kernUpdatePos << > > (numObjects, dt, dev_pos, dev_vel2); + checkCUDAErrorWithLine("kernUpdatePos failed!"); + + std::swap(dev_vel1, dev_vel2); + cudaDeviceSynchronize(); } void Boids::stepSimulationScatteredGrid(float dt) { @@ -374,6 +640,42 @@ void Boids::stepSimulationScatteredGrid(float dt) { // - Perform velocity updates using neighbor search // - Update positions // - Ping-pong buffers as needed + + dim3 fullBlocksPerGrid((numObjects + blockSize - 1) / blockSize); + dim3 blocksCells((gridCellCount + blockSize - 1) / blockSize); + + + kernComputeIndices <<< fullBlocksPerGrid, blockSize >>> ( + numObjects, gridSideCount, gridMinimum, gridInverseCellWidth, + dev_pos, dev_particleArrayIndices, dev_particleGridIndices); + checkCUDAErrorWithLine("kernComputeIndices failed"); + + thrust::sort_by_key( + dev_thrust_particleGridIndices, + dev_thrust_particleGridIndices + numObjects, + dev_thrust_particleArrayIndices); + + kernResetIntBuffer << > > (gridCellCount, dev_gridCellStartIndices, -1); + kernResetIntBuffer << > > (gridCellCount, dev_gridCellEndIndices, -1); + + kernIdentifyCellStartEnd << > > ( + numObjects, dev_particleGridIndices, + dev_gridCellStartIndices, dev_gridCellEndIndices); + checkCUDAErrorWithLine("kernIdentifyCellStartEnd failed"); + + kernUpdateVelNeighborSearchScattered << > > ( + numObjects, gridSideCount, gridMinimum, gridInverseCellWidth, gridCellWidth, + dev_gridCellStartIndices, dev_gridCellEndIndices, + dev_particleArrayIndices, + dev_pos, dev_vel1, dev_vel2); + checkCUDAErrorWithLine("kernUpdateVelNeighborSearchScattered failed"); + + kernUpdatePos << > > (numObjects, dt, dev_pos, dev_vel2); + checkCUDAErrorWithLine("kernUpdatePos failed"); + + std::swap(dev_vel1, dev_vel2); + cudaDeviceSynchronize(); + } void Boids::stepSimulationCoherentGrid(float dt) { @@ -392,14 +694,72 @@ void Boids::stepSimulationCoherentGrid(float dt) { // - Perform velocity updates using neighbor search // - Update positions // - Ping-pong buffers as needed. THIS MAY BE DIFFERENT FROM BEFORE. + dim3 fullBlocksPerGrid((numObjects + blockSize - 1) / blockSize); + dim3 blocksCells((gridCellCount + blockSize - 1) / blockSize); + + kernComputeIndices << > > ( + numObjects, gridSideCount, gridMinimum, gridInverseCellWidth, + dev_pos, dev_particleArrayIndices, dev_particleGridIndices); + checkCUDAErrorWithLine("kernComputeIndices failed"); + + thrust::sort_by_key(dev_thrust_particleGridIndices, + dev_thrust_particleGridIndices + numObjects, + dev_thrust_particleArrayIndices); + + kernReorderDataCoherent << > > ( + numObjects, + dev_particleArrayIndices, + dev_pos, + dev_vel1, + dev_posCoherent, + dev_vel1Coherent); + checkCUDAErrorWithLine("kernReorderDataCoherent failed"); + + kernResetIntBuffer << > > (gridCellCount, dev_gridCellStartIndices, -1); + kernResetIntBuffer << > > (gridCellCount, dev_gridCellEndIndices, -1); + + kernIdentifyCellStartEnd << > > ( + numObjects, dev_particleGridIndices, + dev_gridCellStartIndices, dev_gridCellEndIndices); + checkCUDAErrorWithLine("kernIdentifyCellStartEnd failed"); + + kernUpdateVelNeighborSearchCoherent << > > ( + numObjects, gridSideCount, gridMinimum, gridInverseCellWidth, gridCellWidth, + dev_gridCellStartIndices, dev_gridCellEndIndices, + dev_posCoherent, dev_vel1Coherent, dev_vel2Coherent); + checkCUDAErrorWithLine("kernUpdateVelNeighborSearchCoherent failed"); + + kernUpdatePos << > > (numObjects, dt, dev_posCoherent, dev_vel2Coherent); + checkCUDAErrorWithLine("kernUpdatePos (coherent) failed"); + + std::swap(dev_vel1Coherent, dev_vel2Coherent); + + std::swap(dev_pos, dev_posCoherent); + std::swap(dev_vel1, dev_vel1Coherent); + std::swap(dev_vel2, dev_vel2Coherent); + + cudaDeviceSynchronize(); } +int Boids::getBlockSize() { return threadsPerBlock.x; } +int Boids::getNeighborMode() { return NEIGHBOR_MODE; } +float Boids::getGridCellWidth() { return gridCellWidth; } + + void Boids::endSimulation() { cudaFree(dev_vel1); cudaFree(dev_vel2); 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_posCoherent); + cudaFree(dev_vel1Coherent); + cudaFree(dev_vel2Coherent); } void Boids::unitTest() { diff --git a/src/kernel.h b/src/kernel.h index a38b64d..1be2d1f 100644 --- a/src/kernel.h +++ b/src/kernel.h @@ -9,4 +9,9 @@ namespace Boids { void endSimulation(); void unitTest(); + int getThreadsPerBlock(); + int getBlockSize(); + int getNeighborMode(); + float getGridCellWidth(); + } diff --git a/src/main.cpp b/src/main.cpp index 9c917c0..39b747b 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -17,17 +17,33 @@ #include #include + + // ================ // Configuration // ================ // 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 + + + +#define USE_CUDA_EVENTS 1 + +#if UNIFORM_GRID && COHERENT_GRID +static const char* MODE_STR = "coherent"; +#elif UNIFORM_GRID +static const char* MODE_STR = "scattered"; +#else +static const char* MODE_STR = "naive"; +#endif + + // LOOK-1.2 - change this to adjust particle count in the simulation -const int N_FOR_VIS = 5000; +const int N_FOR_VIS = 10000; const float DT = 0.2f; /** @@ -98,6 +114,7 @@ bool init(int argc, char **argv) { return false; } glfwMakeContextCurrent(window); + glfwSwapInterval(0); glfwSetKeyCallback(window, keyCallback); glfwSetCursorPosCallback(window, mousePositionCallback); glfwSetMouseButtonCallback(window, mouseButtonCallback); @@ -192,33 +209,91 @@ void initShaders(GLuint * program) { //==================================== // Main loop //==================================== + //void runCUDA() { + // // Map OpenGL buffer object for writing from CUDA on a single GPU + // // No data is moved (Win & Linux). When mapped to CUDA, OpenGL should not + // // use this buffer + + // float4 *dptr = NULL; + // float *dptrVertPositions = NULL; + // float *dptrVertVelocities = NULL; + + // cudaGLMapBufferObject((void**)&dptrVertPositions, boidVBO_positions); + // cudaGLMapBufferObject((void**)&dptrVertVelocities, boidVBO_velocities); + + // // execute the kernel + // #if UNIFORM_GRID && COHERENT_GRID + // Boids::stepSimulationCoherentGrid(DT); + // #elif UNIFORM_GRID + // Boids::stepSimulationScatteredGrid(DT); + // #else + // Boids::stepSimulationNaive(DT); + // #endif + + // #if VISUALIZE + // Boids::copyBoidsToVBO(dptrVertPositions, dptrVertVelocities); + // #endif + // // unmap buffer object + // cudaGLUnmapBufferObject(boidVBO_positions); + // cudaGLUnmapBufferObject(boidVBO_velocities); + //} + void runCUDA() { - // Map OpenGL buffer object for writing from CUDA on a single GPU - // No data is moved (Win & Linux). When mapped to CUDA, OpenGL should not - // use this buffer - - float4 *dptr = NULL; - float *dptrVertPositions = NULL; - float *dptrVertVelocities = NULL; - - cudaGLMapBufferObject((void**)&dptrVertPositions, boidVBO_positions); - cudaGLMapBufferObject((void**)&dptrVertVelocities, boidVBO_velocities); - - // execute the kernel - #if UNIFORM_GRID && COHERENT_GRID - Boids::stepSimulationCoherentGrid(DT); - #elif UNIFORM_GRID - Boids::stepSimulationScatteredGrid(DT); - #else - Boids::stepSimulationNaive(DT); - #endif - - #if VISUALIZE - Boids::copyBoidsToVBO(dptrVertPositions, dptrVertVelocities); - #endif - // unmap buffer object - cudaGLUnmapBufferObject(boidVBO_positions); - cudaGLUnmapBufferObject(boidVBO_velocities); + +#if VISUALIZE + float* dptrVertPositions = nullptr; + float* dptrVertVelocities = nullptr; + cudaGLMapBufferObject((void**)&dptrVertPositions, boidVBO_positions); + cudaGLMapBufferObject((void**)&dptrVertVelocities, boidVBO_velocities); +#endif + +#if USE_CUDA_EVENTS + static bool evInit = false; + static cudaEvent_t evStart, evStop; + static float acc_ms = 0.0f; + static int acc_n = 0; + if (!evInit) { cudaEventCreate(&evStart); cudaEventCreate(&evStop); evInit = true; } + cudaEventRecord(evStart); +#endif + +#if UNIFORM_GRID && COHERENT_GRID + Boids::stepSimulationCoherentGrid(DT); +#elif UNIFORM_GRID + Boids::stepSimulationScatteredGrid(DT); +#else + Boids::stepSimulationNaive(DT); +#endif + +#if USE_CUDA_EVENTS + cudaEventRecord(evStop); + cudaEventSynchronize(evStop); + float ms = 0.f; + cudaEventElapsedTime(&ms, evStart, evStop); + acc_ms += ms; + acc_n += 1; + + + if (acc_n == 120) { + float avg = acc_ms / acc_n; + printf("mode=%s,vis=%d,N=%d,block=%d,nb=%d,cellw=%.2f,ms=%.6f,fps=%.2f\n", + MODE_STR, (int)VISUALIZE, N_FOR_VIS, + Boids::getBlockSize(), + Boids::getNeighborMode(), + Boids::getGridCellWidth(), + avg, 1000.0 / avg); + fflush(stdout); + + acc_ms = 0.f; + acc_n = 0; + } +#endif + +#if VISUALIZE + + Boids::copyBoidsToVBO(dptrVertPositions, dptrVertVelocities); + cudaGLUnmapBufferObject(boidVBO_positions); + cudaGLUnmapBufferObject(boidVBO_velocities); +#endif } void mainLoop() {