diff --git a/GRAPHING.md b/GRAPHING.md new file mode 100644 index 0000000..771ec93 --- /dev/null +++ b/GRAPHING.md @@ -0,0 +1,120 @@ +GRAPHING WITH PYTHON +==================== + +## Files + +The python script may be found [here](outputData/csvReader.py). The other relevant code is in the `main` file within `src`. For any gaps, or if you're curious as to the support structure behind some of this functionality, feel free to look there. + +## Recording the CSV file + +After collecting the data across a number of iterations, inside a `std::vector` of structs called `eventRecords`, I looped through each of them to extract the iteration number, the time between the two records, and the total simulation time up until that point: + +```C++ + void writeTime(const char* fileName) { + FILE* of = fopen(fileName, "w"); + + for (auto& record : eventRecords) { + double millisPerFrame = record.time / TIMEKEEPING_FRAMESIZE; + millisPerFrame /= 1000.0;//seconds per frame + double fps = 1.0 / millisPerFrame; + double seconds = record.totalTime / 1000.0; + fprintf(of, "%d,%0.3f,%f\n", record.frameNo, seconds, fps); + }//for + + fclose(of); + }//writeTime +``` + +After this, I wrote the records into the csv. This involves rows being separated by newlines, and columns being separated by commas. + +For my particular code, the filenames were hard-coded filenames, like the following: + +```C++ + const char timingFileName[] = "../outputData/CoherentGrid_HighDensity_128.csv"; +``` + +## Using the Data + +### Invoking the Python script + +The script I wrote takes in at least two arguments; the first is the title for the graph (and, by extension, the name of the image to save), and all arguments after that are names of csv files, in the format recorded above. + +For example, a way to use this script would be to call `python csvReader.py "Coherent Grid, High Density, All Block Sizes" CoherentGrid_HighDensity_*.csv` + +The main function is as follows: + +```Python + + def main(): + if len(sys.argv) < 3: + print("Please input a title and file names") + exit(0) + + resultSets = [] + + for i, fileName in enumerate(sys.argv): + if i == 0 or i == 1: + continue + resultSets.append((fileName, readCSV(fileName))) + + makeGraphs(resultSets, sys.argv[1]) + +``` + +It takes each CSV file, transforms it into a Numpy array, and binds that into a tuple along with the name of the file the data came from. It then hands each of those sets off to a function that graphs them, and saves out the data (or, optionally, displays it onscreen). + +### Reading in the CSV data + +I made use of the `csv` package, but the file format is so simple, you could quickly write up your own. The function was the following: + +```Python + def readCSV(filename): + """ + takes in a filename + returns a numpy array, with the first row as the + timestamp in seconds, and the second row + as the fps across the last time block + """ + results = [] + with open(filename) as csv_file: + reader = csv.reader(csv_file, delimiter=',') + for line in reader: + timestamp = float(line[1]) + fps = float(line[2]) + results.append([timestamp, fps]) + + return np.array(results).T +``` + +If you're not familiar with the `numpy` library, it's a handy library for making array operations more efficient, while also making some simple things entirely too complicated. In this case, the `.T` within the return statement transposes the 2D array, making our two columns into two rows, which will be essential for displaying them nicely. + +### Graphing the data + +I used `matplotlib` to graph the data. It has a billion fun features, along with occasionally frustrating documentation, meaning all of my matplotlib code is a hacked-together mess of code snippets stolen from forum posts. That said, the goals here were relatively simple; I wanted to take a bunch of 2d arrays representing pairs of `(timestamp, fps)` data points and graph those series' onto a line plot. + +The code was as follows: + +```Python + def makeGraphs(resultSets, title): + """ + Displays the resultant data sets, along with a given title + """ + fig, ax = plt.subplots(1) + for filename, data in resultSets: + ax.plot(data[0], data[1], label = cleanFileName(filename)) + + ax.legend() + plt.xlabel("Time (seconds)") + plt.ylabel("Ticks/Frames per Second") + + fig.suptitle(title) + fig.set_size_inches(10,6) + + #plt.show() #uncomment this to display the graph on your screen + filePath = makeSavePath(title) + plt.savefig(filePath) +``` + +The core functionality is the `plot` function, where you can provide a series of `x` data, along with a series of `y` data, along with a bunch of other optional variables, such as a label for the series. I did this for each set of data, stuck some labels onto the plot, and then saved it all to an image file. + +Happy coding! diff --git a/README.md b/README.md index d63a6a1..65d389b 100644 --- a/README.md +++ b/README.md @@ -1,11 +1,67 @@ **University of Pennsylvania, CIS 565: 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) +* Taylor Nelms + * [LinkedIn](https://www.linkedin.com/in/taylor-k-7b2110191/), [twitter](https://twitter.com/nelms_taylor) +* Tested on: Windows 10, Intel i3 Coffee Lake 4-core 3.6GHz processor, 16GB RAM, NVidia GeForce GTX1650 4GB -### (TODO: Your README) +## Results +![](images/mLo_dMed_gMed.gif) +*No Grid used for implementation* +![](images/mMed_dMed_gMed.gif) +*Uniform grid used within implementation* +![](images/mHi_dMed_gMed.gif) +*Coherent grid used within implementation* -Include screenshots, analysis, etc. (Remember, this is public, so don't put -anything here that you don't want to share with the world.) +## Analysis + +### Implementation Strategy + +Unsurprisingly, the grid implementations ended up significantly more efficient than the naive implementation. For runs with 5000 boids, with a block size of 128, the FPS over a 45-ish-second run yielded the following results: + +![](images/All Grids, Medium Density, Block Size 128.png) + +There are a few things to unpack here. First, the spike in the initial framerate of the naive implementation. Frankly, I have no idea why this exists; I was taking data points every 300 ticks, so I can't imagine this being some fluke of initial set-up taking less time. In all honesty, I would need to do significantly more debugging to figure it out. + +Of course, the more interesting behavior lies within the meat of the simulation. The grid-based solutions performed better overall, with a slight improvement for the coherent grid over the uniform grid. Of course, algorithmically, this makes sense; the grid-based approaches cut execution time on the GPU from order `O(N)` to `O(n)`, where `n` is the number of boids that are within the grid-neighborhood of each boid. (On a CPU, the naive approach would run in time `O(N^2)`, while the grid approaches would run in time `O(Nn)`.) + +For another example, let's look at each of the models with a higher density of boids; in this case, we're operating with 10,000 boids, rather than 5,000, in that same space: + +![](images/All Grids, High Density, Block Size 128.png) + +Another notable part here is that the framerate drops off over time for the uniform grid, while the coherent grid stays relatively steady. The best I can figure for the drop is that, over time, each boid has more neighbors, and so the number of data accesses to those neighbors increases (as more of the boids are in flocks). This increases the penalties felt from the boid data being more scattered in the uniform grid implementation, as cached data accesses become less favorable. + +### Number of Boids + +Unsurprisingly, as the number of boids increase, the execution speed of the simulation decreases. Here are some comparisons for all the models, running with `2000` boids, `5000` boids, and `10000` boids: + +![](images/No Grid, All Densities, Block Size 128.png) ![](images/Uniform Grid, All Densities, Block Size 128.png) ![](images/Coherent Grid, All Densities, Block Size 128.png) + +Unsurprisingly, the naive implementation shows a roughly linear relationship between simulation speed and the number of boids. The others have a more complex relationship, but the overall trend is clear, and they seem to do slightly better than linear with sample size. + +### Block Size + +The differences between block size were very interesting. I ran a series of simulations with block sizes of `32`, `128`, and `512`. Here are a couple of graphs comparing runs with various block sizes: + +![](images/No Grid, Medium Density, All Block Sizes.png)![](images/Coherent Grid, Medium Density, All Block Sizes.png) + +Notably, there is not that much difference in performance for the naive implementation based on block size. This makes some sense; so many of the operations are simple and easily parallelizable, it is hard to imagine that the various levels of scheduling or memory caching could make a significant performance difference. + +However, for the grid implementations, block sizes made huge differences in outcome. + +The block size of 32 ran the worst. This makes some sense; there must be some amount of overhead in creating a block, and getting access to the relevant memory, and given the number of blocks that were spun up at various points in the simulation step, it makes sense that those penalties would add up. + +This would imply that a larger block size would improve performance; however, we see performance dip when we increase the block size from `128` to `512`. The best explanation I can think of is that a block needing to finish together before a new block can be put in its place would lead to situations where an entire block could be held up by a few rogure warps. In those cases, a whole section of processing power could be lost while the scheduler keeps a block running. + +### Bonus Graph + +Everyone needs a little graph gore in their life every now and then: + +![](images/All Test Runs.png) + +### Miscellaneous Notes + +During the naive implementation, I changed my `distBetween` function, which computed the distance between two vectors, +between using the `glm::distance` function and a simple `sqrt(xdiff * xdiff + ydiff * ydiff + zdiff * zdiff)` function. +Though I would have expected the `glm::distance` function to be highly optimized in some fashion, +I saw framerate drop from around 10fps to around 2.5fps in the simulation window. diff --git a/images/All Grids, High Density, Block Size 128.png b/images/All Grids, High Density, Block Size 128.png new file mode 100644 index 0000000..fd0b251 Binary files /dev/null and b/images/All Grids, High Density, Block Size 128.png differ diff --git a/images/All Grids, Medium Density, Block Size 128.png b/images/All Grids, Medium Density, Block Size 128.png new file mode 100644 index 0000000..24a4fe0 Binary files /dev/null and b/images/All Grids, Medium Density, Block Size 128.png differ diff --git a/images/All Test Runs.png b/images/All Test Runs.png new file mode 100644 index 0000000..5a80d95 Binary files /dev/null and b/images/All Test Runs.png differ diff --git a/images/Coherent Grid, All Densities, All Block Sizes.png b/images/Coherent Grid, All Densities, All Block Sizes.png new file mode 100644 index 0000000..0e8185b Binary files /dev/null and b/images/Coherent Grid, All Densities, All Block Sizes.png differ diff --git a/images/Coherent Grid, All Densities, Block Size 128.png b/images/Coherent Grid, All Densities, Block Size 128.png new file mode 100644 index 0000000..36db4eb Binary files /dev/null and b/images/Coherent Grid, All Densities, Block Size 128.png differ diff --git a/images/Coherent Grid, High Density, All Block Sizes.png b/images/Coherent Grid, High Density, All Block Sizes.png new file mode 100644 index 0000000..a51c391 Binary files /dev/null and b/images/Coherent Grid, High Density, All Block Sizes.png differ diff --git a/images/Coherent Grid, Medium Density, All Block Sizes.png b/images/Coherent Grid, Medium Density, All Block Sizes.png new file mode 100644 index 0000000..2f42580 Binary files /dev/null and b/images/Coherent Grid, Medium Density, All Block Sizes.png differ diff --git a/images/No Grid, All Densities, All Block Sizes.png b/images/No Grid, All Densities, All Block Sizes.png new file mode 100644 index 0000000..fd3afe7 Binary files /dev/null and b/images/No Grid, All Densities, All Block Sizes.png differ diff --git a/images/No Grid, All Densities, Block Size 128.png b/images/No Grid, All Densities, Block Size 128.png new file mode 100644 index 0000000..ff324eb Binary files /dev/null and b/images/No Grid, All Densities, Block Size 128.png differ diff --git a/images/No Grid, Medium Density, All Block Sizes.png b/images/No Grid, Medium Density, All Block Sizes.png new file mode 100644 index 0000000..2819bfe Binary files /dev/null and b/images/No Grid, Medium Density, All Block Sizes.png differ diff --git a/images/Uniform Grid, All Densities, All Block Sizes.png b/images/Uniform Grid, All Densities, All Block Sizes.png new file mode 100644 index 0000000..28d7ec4 Binary files /dev/null and b/images/Uniform Grid, All Densities, All Block Sizes.png differ diff --git a/images/Uniform Grid, All Densities, Block Size 128.png b/images/Uniform Grid, All Densities, Block Size 128.png new file mode 100644 index 0000000..2c1864f Binary files /dev/null and b/images/Uniform Grid, All Densities, Block Size 128.png differ diff --git a/images/mHi_dMed_gMed.gif b/images/mHi_dMed_gMed.gif new file mode 100644 index 0000000..38fe153 Binary files /dev/null and b/images/mHi_dMed_gMed.gif differ diff --git a/images/mLo_dMed_gMed.gif b/images/mLo_dMed_gMed.gif new file mode 100644 index 0000000..42a5cde Binary files /dev/null and b/images/mLo_dMed_gMed.gif differ diff --git a/images/mMed_dMed_gMed.gif b/images/mMed_dMed_gMed.gif new file mode 100644 index 0000000..df1b569 Binary files /dev/null and b/images/mMed_dMed_gMed.gif differ diff --git a/outputData/CoherentGrid_HighDensity_128.csv b/outputData/CoherentGrid_HighDensity_128.csv new file mode 100644 index 0000000..dd21c2d --- /dev/null +++ b/outputData/CoherentGrid_HighDensity_128.csv @@ -0,0 +1,69 @@ +300,0.698,429.799427 +600,1.326,477.707006 +900,1.976,462.249615 +1200,2.618,468.750000 +1500,3.268,461.538462 +1800,3.912,466.562986 +2100,4.568,458.015267 +2400,5.214,465.116279 +2700,5.870,458.015267 +3000,6.519,462.962963 +3300,7.193,445.765230 +3600,7.853,455.235205 +3900,8.510,457.317073 +4200,9.166,458.015267 +4500,9.826,455.235205 +4800,10.487,454.545455 +5100,11.161,445.765230 +5400,11.841,441.826215 +5700,12.503,453.857791 +6000,13.168,451.807229 +6300,13.814,465.116279 +6600,14.463,462.962963 +6900,15.133,449.101796 +7200,15.802,449.101796 +7500,16.460,455.927052 +7800,17.124,452.488688 +8100,17.782,456.621005 +8400,18.435,459.418070 +8700,19.082,465.116279 +9000,19.766,439.238653 +9300,20.421,458.715596 +9600,21.074,460.122699 +9900,21.723,462.962963 +10200,22.380,457.317073 +10500,23.052,447.093890 +10800,23.715,453.857791 +11100,24.386,447.761194 +11400,25.038,460.122699 +11700,25.701,452.488688 +12000,26.366,452.488688 +12300,27.036,448.430493 +12600,27.690,459.418070 +12900,28.353,453.172205 +13200,29.018,451.807229 +13500,29.673,458.015267 +13800,30.327,459.418070 +14100,30.987,455.235205 +14400,31.643,458.015267 +14700,32.314,447.093890 +15000,32.973,456.621005 +15300,33.626,460.122699 +15600,34.292,451.127820 +15900,34.958,451.127820 +16200,35.641,439.882698 +16500,36.305,452.488688 +16800,36.975,448.430493 +17100,37.640,451.807229 +17400,38.296,458.715596 +17700,38.944,462.962963 +18000,39.616,447.093890 +18300,40.291,445.765230 +18600,40.941,462.249615 +18900,41.591,461.538462 +19200,42.249,456.621005 +19500,42.924,445.103858 +19800,43.597,447.093890 +20100,44.261,452.488688 +20400,44.928,449.775112 +20700,45.593,452.488688 diff --git a/outputData/CoherentGrid_HighDensity_512.csv b/outputData/CoherentGrid_HighDensity_512.csv new file mode 100644 index 0000000..e6439ce --- /dev/null +++ b/outputData/CoherentGrid_HighDensity_512.csv @@ -0,0 +1,64 @@ +300,0.728,412.087912 +600,1.408,441.826215 +900,2.087,442.477876 +1200,2.779,434.153401 +1500,3.502,414.937759 +1800,4.204,427.960057 +2100,4.907,426.742532 +2400,5.594,437.317784 +2700,6.288,432.900433 +3000,6.988,429.184549 +3300,7.685,430.416069 +3600,8.400,420.168067 +3900,9.119,417.827298 +4200,9.820,428.571429 +4500,10.515,431.654676 +4800,11.220,426.742532 +5100,11.929,423.131171 +5400,12.630,428.571429 +5700,13.330,429.799427 +6000,14.014,438.596491 +6300,14.727,421.940928 +6600,15.431,426.742532 +6900,16.120,436.046512 +7200,16.801,441.176471 +7500,17.476,444.444444 +7800,18.184,424.328147 +8100,18.884,429.184549 +8400,19.583,429.799427 +8700,20.297,420.757363 +9000,21.007,423.131171 +9300,21.713,425.531915 +9600,22.431,418.410042 +9900,23.131,429.184549 +10200,23.840,423.728814 +10500,24.537,431.034483 +10800,25.239,427.960057 +11100,25.958,417.246175 +11400,26.672,421.348315 +11700,27.366,432.900433 +12000,28.066,428.571429 +12300,28.764,430.416069 +12600,29.476,421.940928 +12900,30.146,448.430493 +13200,30.857,422.535211 +13500,31.548,434.153401 +13800,32.258,423.131171 +14100,32.947,436.046512 +14400,33.643,431.654676 +14700,34.342,429.799427 +15000,35.020,442.477876 +15300,35.734,421.348315 +15600,36.460,413.223140 +15900,37.155,432.276657 +16200,37.882,413.223140 +16500,38.586,426.742532 +16800,39.296,423.131171 +17100,40.014,417.827298 +17400,40.721,424.929178 +17700,41.443,416.088766 +18000,42.145,428.571429 +18300,42.837,434.153401 +18600,43.552,420.168067 +18900,44.277,414.937759 +19200,44.986,423.728814 diff --git a/outputData/CoherentGrid_LowDensity_128.csv b/outputData/CoherentGrid_LowDensity_128.csv new file mode 100644 index 0000000..e5aa434 --- /dev/null +++ b/outputData/CoherentGrid_LowDensity_128.csv @@ -0,0 +1,137 @@ +300,0.411,729.927007 +600,0.752,879.765396 +900,1.075,931.677019 +1200,1.409,900.900901 +1500,1.725,952.380952 +1800,2.049,928.792570 +2100,2.379,911.854103 +2400,2.701,934.579439 +2700,3.023,934.579439 +3000,3.352,914.634146 +3300,3.676,928.792570 +3600,3.996,937.500000 +3900,4.318,934.579439 +4200,4.639,937.500000 +4500,4.965,923.076923 +4800,5.289,928.792570 +5100,5.619,909.090909 +5400,5.943,928.792570 +5700,6.270,920.245399 +6000,6.599,911.854103 +6300,6.924,928.792570 +6600,7.250,923.076923 +6900,7.574,928.792570 +7200,7.906,906.344411 +7500,8.237,906.344411 +7800,8.560,931.677019 +8100,8.880,940.438871 +8400,9.217,890.207715 +8700,9.541,928.792570 +9000,9.870,917.431193 +9300,10.192,931.677019 +9600,10.522,911.854103 +9900,10.850,917.431193 +10200,11.180,911.854103 +10500,11.509,911.854103 +10800,11.841,909.090909 +11100,12.171,911.854103 +11400,12.496,923.076923 +11700,12.825,914.634146 +12000,13.158,906.344411 +12300,13.504,867.052023 +12600,13.833,914.634146 +12900,14.160,920.245399 +13200,14.489,917.431193 +13500,14.817,917.431193 +13800,15.145,914.634146 +14100,15.472,920.245399 +14400,15.799,920.245399 +14700,16.128,914.634146 +15000,16.459,909.090909 +15300,16.797,887.573964 +15600,17.124,923.076923 +15900,17.458,900.900901 +16200,17.793,898.203593 +16500,18.121,917.431193 +16800,18.446,923.076923 +17100,18.779,903.614458 +17400,19.110,909.090909 +17700,19.440,911.854103 +18000,19.766,923.076923 +18300,20.094,917.431193 +18600,20.420,923.076923 +18900,20.764,872.093023 +19200,21.093,914.634146 +19500,21.419,920.245399 +19800,21.748,914.634146 +20100,22.070,931.677019 +20400,22.405,898.203593 +20700,22.733,920.245399 +21000,23.060,917.431193 +21300,23.385,925.925926 +21600,23.718,903.614458 +21900,24.046,917.431193 +22200,24.382,895.522388 +22500,24.720,890.207715 +22800,25.052,906.344411 +23100,25.385,903.614458 +23400,25.707,934.579439 +23700,26.044,892.857143 +24000,26.375,909.090909 +24300,26.708,900.900901 +24600,27.031,931.677019 +24900,27.355,925.925926 +25200,27.691,895.522388 +25500,28.019,914.634146 +25800,28.352,903.614458 +26100,28.680,917.431193 +26400,29.013,903.614458 +26700,29.351,887.573964 +27000,29.688,892.857143 +27300,30.007,940.438871 +27600,30.338,909.090909 +27900,30.671,903.614458 +28200,31.004,903.614458 +28500,31.349,872.093023 +28800,31.709,835.654596 +29100,32.045,895.522388 +29400,32.373,917.431193 +29700,32.707,900.900901 +30000,33.032,925.925926 +30300,33.363,909.090909 +30600,33.701,890.207715 +30900,34.034,903.614458 +31200,34.370,895.522388 +31500,34.704,900.900901 +31800,35.037,903.614458 +32100,35.376,887.573964 +32400,35.707,909.090909 +32700,36.032,923.076923 +33000,36.365,906.344411 +33300,36.694,914.634146 +33600,37.030,895.522388 +33900,37.358,917.431193 +34200,37.687,911.854103 +34500,38.033,869.565217 +34800,38.362,911.854103 +35100,38.705,877.192982 +35400,39.036,909.090909 +35700,39.371,898.203593 +36000,39.709,890.207715 +36300,40.038,914.634146 +36600,40.367,911.854103 +36900,40.696,917.431193 +37200,41.033,890.207715 +37500,41.366,903.614458 +37800,41.691,925.925926 +38100,42.026,898.203593 +38400,42.356,911.854103 +38700,42.693,892.857143 +39000,43.025,906.344411 +39300,43.348,931.677019 +39600,43.691,877.192982 +39900,44.022,906.344411 +40200,44.365,877.192982 +40500,44.700,898.203593 +40800,45.036,895.522388 +41100,45.375,887.573964 diff --git a/outputData/CoherentGrid_MedDensity_128.csv b/outputData/CoherentGrid_MedDensity_128.csv new file mode 100644 index 0000000..a1ecb94 --- /dev/null +++ b/outputData/CoherentGrid_MedDensity_128.csv @@ -0,0 +1,102 @@ +300,0.518,579.150579 +600,0.959,681.818182 +900,1.402,680.272109 +1200,1.837,689.655172 +1500,2.284,672.645740 +1800,2.722,686.498856 +2100,3.158,688.073394 +2400,3.604,674.157303 +2700,4.050,674.157303 +3000,4.505,660.792952 +3300,4.943,686.498856 +3600,5.387,677.200903 +3900,5.834,674.157303 +4200,6.279,674.157303 +4500,6.728,669.642857 +4800,7.164,689.655172 +5100,7.606,680.272109 +5400,8.047,681.818182 +5700,8.493,674.157303 +6000,8.936,678.733032 +6300,9.384,671.140940 +6600,9.831,672.645740 +6900,10.273,680.272109 +7200,10.719,674.157303 +7500,11.163,677.200903 +7800,11.613,668.151448 +8100,12.059,675.675676 +8400,12.497,684.931507 +8700,12.948,666.666667 +9000,13.404,659.340659 +9300,13.859,660.792952 +9600,14.310,665.188470 +9900,14.766,659.340659 +10200,15.218,665.188470 +10500,15.674,659.340659 +10800,16.128,660.792952 +11100,16.573,675.675676 +11400,17.022,669.642857 +11700,17.468,672.645740 +12000,17.918,666.666667 +12300,18.368,668.151448 +12600,18.821,663.716814 +12900,19.269,671.140940 +13200,19.723,660.792952 +13500,20.171,671.140940 +13800,20.627,659.340659 +14100,21.074,671.140940 +14400,21.522,671.140940 +14700,21.972,668.151448 +15000,22.428,657.894737 +15300,22.878,668.151448 +15600,23.333,659.340659 +15900,23.783,668.151448 +16200,24.229,674.157303 +16500,24.675,674.157303 +16800,25.125,666.666667 +17100,25.576,666.666667 +17400,26.029,663.716814 +17700,26.476,672.645740 +18000,26.924,671.140940 +18300,27.376,665.188470 +18600,27.827,666.666667 +18900,28.288,652.173913 +19200,28.741,663.716814 +19500,29.196,660.792952 +19800,29.656,653.594771 +20100,30.108,663.716814 +20400,30.559,666.666667 +20700,31.007,671.140940 +21000,31.455,671.140940 +21300,31.930,632.911392 +21600,32.389,655.021834 +21900,32.845,659.340659 +22200,33.297,663.716814 +22500,33.746,669.642857 +22800,34.198,666.666667 +23100,34.654,659.340659 +23400,35.098,677.200903 +23700,35.545,672.645740 +24000,35.987,678.733032 +24300,36.435,671.140940 +24600,36.886,666.666667 +24900,37.336,668.151448 +25200,37.788,665.188470 +25500,38.232,675.675676 +25800,38.679,674.157303 +26100,39.119,681.818182 +26400,39.567,671.140940 +26700,40.018,666.666667 +27000,40.454,689.655172 +27300,40.907,663.716814 +27600,41.357,666.666667 +27900,41.810,663.716814 +28200,42.262,665.188470 +28500,42.722,653.594771 +28800,43.172,668.151448 +29100,43.622,668.151448 +29400,44.068,674.157303 +29700,44.511,677.200903 +30000,44.967,659.340659 +30300,45.421,662.251656 +30600,45.884,649.350649 diff --git a/outputData/CoherentGrid_MedDensity_32.csv b/outputData/CoherentGrid_MedDensity_32.csv new file mode 100644 index 0000000..f435af5 --- /dev/null +++ b/outputData/CoherentGrid_MedDensity_32.csv @@ -0,0 +1,76 @@ +300,0.647,463.678516 +600,1.227,518.134715 +900,1.821,505.902192 +1200,2.396,522.648084 +1500,2.998,499.168053 +1800,3.594,504.201681 +2100,4.202,493.421053 +2400,4.794,507.614213 +2700,5.392,502.512563 +3000,6.004,490.998363 +3300,6.592,511.073254 +3600,7.188,503.355705 +3900,7.785,503.355705 +4200,8.385,500.834725 +4500,9.000,488.599349 +4800,9.603,499.168053 +5100,10.198,505.050505 +5400,10.796,502.512563 +5700,11.390,505.050505 +6000,12.000,492.610837 +6300,12.584,514.579760 +6600,13.196,490.998363 +6900,13.807,492.610837 +7200,14.412,496.688742 +7500,15.006,505.902192 +7800,15.598,507.614213 +8100,16.189,507.614213 +8400,16.791,500.000000 +8700,17.393,499.168053 +9000,17.988,505.050505 +9300,18.588,500.834725 +9600,19.194,495.867769 +9900,19.801,495.049505 +10200,20.389,511.073254 +10500,20.997,494.233937 +10800,21.598,499.168053 +11100,22.197,501.672241 +11400,22.786,511.073254 +11700,23.378,507.614213 +12000,23.975,503.355705 +12300,24.571,504.201681 +12600,25.174,498.338870 +12900,25.785,491.803279 +13200,26.398,489.396411 +13500,27.009,491.803279 +13800,27.618,494.233937 +14100,28.204,512.820513 +14400,28.812,494.233937 +14700,29.444,475.435816 +15000,30.047,499.168053 +15300,30.655,494.233937 +15600,31.248,505.902192 +15900,31.858,492.610837 +16200,32.456,502.512563 +16500,33.059,498.338870 +16800,33.661,499.168053 +17100,34.249,511.073254 +17400,34.860,491.803279 +17700,35.462,499.168053 +18000,36.071,493.421053 +18300,36.674,497.512438 +18600,37.268,505.902192 +18900,37.874,495.867769 +19200,38.474,500.834725 +19500,39.069,505.050505 +19800,39.664,505.050505 +20100,40.262,501.672241 +20400,40.867,496.688742 +20700,41.478,491.803279 +21000,42.087,493.421053 +21300,42.680,506.756757 +21600,43.298,486.223663 +21900,43.915,487.012987 +22200,44.515,500.834725 +22500,45.131,487.804878 +22800,45.736,496.688742 diff --git a/outputData/CoherentGrid_MedDensity_512.csv b/outputData/CoherentGrid_MedDensity_512.csv new file mode 100644 index 0000000..0d71c8d --- /dev/null +++ b/outputData/CoherentGrid_MedDensity_512.csv @@ -0,0 +1,95 @@ +300,0.545,550.458716 +600,1.015,639.658849 +900,1.479,647.948164 +1200,1.945,643.776824 +1500,2.412,643.776824 +1800,2.878,645.161290 +2100,3.351,635.593220 +2400,3.818,645.161290 +2700,4.282,646.551724 +3000,4.755,634.249471 +3300,5.222,645.161290 +3600,5.691,641.025641 +3900,6.162,638.297872 +4200,6.638,631.578947 +4500,7.111,635.593220 +4800,7.577,645.161290 +5100,8.050,634.249471 +5400,8.527,630.252101 +5700,8.990,649.350649 +6000,9.453,647.948164 +6300,9.920,645.161290 +6600,10.398,628.930818 +6900,10.872,634.249471 +7200,11.331,655.021834 +7500,11.801,639.658849 +7800,12.276,632.911392 +8100,12.733,657.894737 +8400,13.194,652.173913 +8700,13.658,647.948164 +9000,14.132,635.593220 +9300,14.605,634.249471 +9600,15.076,638.297872 +9900,15.537,652.173913 +10200,16.021,621.118012 +10500,16.487,645.161290 +10800,16.954,643.776824 +11100,17.423,641.025641 +11400,17.881,656.455142 +11700,18.365,621.118012 +12000,18.838,635.593220 +12300,19.312,634.249471 +12600,19.783,638.297872 +12900,20.253,639.658849 +13200,20.729,631.578947 +13500,21.197,642.398287 +13800,21.652,660.792952 +14100,22.127,631.578947 +14400,22.596,641.025641 +14700,23.058,650.759219 +15000,23.517,653.594771 +15300,23.978,652.173913 +15600,24.450,638.297872 +15900,24.915,646.551724 +16200,25.383,642.398287 +16500,25.853,639.658849 +16800,26.322,641.025641 +17100,26.790,642.398287 +17400,27.253,650.759219 +17700,27.739,617.283951 +18000,28.224,619.834711 +18300,28.699,631.578947 +18600,29.170,638.297872 +18900,29.661,612.244898 +19200,30.131,639.658849 +19500,30.615,621.118012 +19800,31.088,635.593220 +20100,31.556,642.398287 +20400,32.042,619.834711 +20700,32.501,655.021834 +21000,32.963,650.759219 +21300,33.441,628.930818 +21600,33.915,634.249471 +21900,34.377,649.350649 +22200,34.851,634.249471 +22500,35.317,645.161290 +22800,35.779,650.759219 +23100,36.247,642.398287 +23400,36.729,623.700624 +23700,37.194,646.551724 +24000,37.671,630.252101 +24300,38.146,631.578947 +24600,38.623,631.578947 +24900,39.092,639.658849 +25200,39.570,628.930818 +25500,40.038,642.398287 +25800,40.512,634.249471 +26100,40.988,631.578947 +26400,41.450,649.350649 +26700,41.934,621.118012 +27000,42.411,630.252101 +27300,42.886,631.578947 +27600,43.372,618.556701 +27900,43.854,623.700624 +28200,44.347,609.756098 +28500,44.823,631.578947 diff --git a/outputData/NoGrid_HighDensity_128.csv b/outputData/NoGrid_HighDensity_128.csv new file mode 100644 index 0000000..44e5da1 --- /dev/null +++ b/outputData/NoGrid_HighDensity_128.csv @@ -0,0 +1,28 @@ +300,0.218,1376.146789 +600,0.544,923.076923 +900,2.245,176.678445 +1200,3.980,174.216028 +1500,5.695,176.366843 +1800,7.394,176.678445 +2100,9.094,177.514793 +2400,10.809,176.470588 +2700,12.494,178.253119 +3000,14.208,176.470588 +3300,15.919,175.438596 +3600,17.623,177.725118 +3900,19.310,178.041543 +4200,21.020,176.991150 +4500,22.706,178.041543 +4800,24.420,176.159718 +5100,26.122,176.470588 +5400,27.842,175.849941 +5700,29.538,176.886792 +6000,31.231,177.304965 +6300,32.949,174.723355 +6600,34.665,176.056338 +6900,36.360,177.200236 +7200,38.059,178.041543 +7500,39.774,175.029172 +7800,41.472,176.678445 +8100,43.173,178.359096 +8400,44.897,175.541252 diff --git a/outputData/NoGrid_HighDensity_512.csv b/outputData/NoGrid_HighDensity_512.csv new file mode 100644 index 0000000..b1b94bc --- /dev/null +++ b/outputData/NoGrid_HighDensity_512.csv @@ -0,0 +1,25 @@ +300,0.223,1345.291480 +600,0.608,779.220779 +900,2.581,153.139357 +1200,4.554,153.217569 +1500,6.520,152.671756 +1800,8.485,152.749491 +2100,10.462,151.975684 +2400,12.428,152.671756 +2700,14.393,153.846154 +3000,16.357,152.905199 +3300,18.321,152.827305 +3600,20.287,152.749491 +3900,22.226,154.958678 +4200,24.200,153.295861 +4500,26.154,153.767299 +4800,28.123,153.295861 +5100,30.097,153.139357 +5400,32.056,153.217569 +5700,34.017,153.139357 +6000,35.986,153.688525 +6300,37.934,154.162384 +6600,39.902,152.516523 +6900,41.868,153.531218 +7200,43.824,153.452685 +7500,45.794,152.361605 diff --git a/outputData/NoGrid_LowDensity_128.csv b/outputData/NoGrid_LowDensity_128.csv new file mode 100644 index 0000000..1676ea9 --- /dev/null +++ b/outputData/NoGrid_LowDensity_128.csv @@ -0,0 +1,138 @@ +300,0.181,1657.458564 +600,0.337,1935.483871 +900,0.460,2459.016393 +1200,0.766,983.606557 +1500,1.091,925.925926 +1800,1.424,900.900901 +2100,1.767,879.765396 +2400,2.105,890.207715 +2700,2.449,877.192982 +3000,2.785,892.857143 +3300,3.115,914.634146 +3600,3.468,852.272727 +3900,3.793,925.925926 +4200,4.128,900.900901 +4500,4.470,877.192982 +4800,4.810,887.573964 +5100,5.125,955.414013 +5400,5.453,917.431193 +5700,5.800,867.052023 +6000,6.145,874.635569 +6300,6.485,882.352941 +6600,6.820,895.522388 +6900,7.173,852.272727 +7200,7.495,937.500000 +7500,7.848,852.272727 +7800,8.172,928.792570 +8100,8.504,914.634146 +8400,8.852,862.068966 +8700,9.183,909.090909 +9000,9.531,862.068966 +9300,9.866,898.203593 +9600,10.202,898.203593 +9900,10.545,874.635569 +10200,10.894,862.068966 +10500,11.237,877.192982 +10800,11.568,909.090909 +11100,11.914,869.565217 +11400,12.256,882.352941 +11700,12.593,895.522388 +12000,12.946,854.700855 +12300,13.261,958.466454 +12600,13.596,895.522388 +12900,13.946,862.068966 +13200,14.297,854.700855 +13500,14.624,923.076923 +13800,14.971,869.565217 +14100,15.307,895.522388 +14400,15.641,900.900901 +14700,15.985,877.192982 +15000,16.313,917.431193 +15300,16.646,906.344411 +15600,16.959,961.538462 +15900,17.307,864.553314 +16200,17.632,925.925926 +16500,17.973,882.352941 +16800,18.335,847.457627 +17100,18.659,931.677019 +17400,19.008,887.573964 +17700,19.361,852.272727 +18000,19.707,872.093023 +18300,20.038,909.090909 +18600,20.377,887.573964 +18900,20.721,874.635569 +19200,21.066,869.565217 +19500,21.402,898.203593 +19800,21.741,887.573964 +20100,22.085,877.192982 +20400,22.416,909.090909 +20700,22.761,874.635569 +21000,23.101,884.955752 +21300,23.448,867.052023 +21600,23.776,920.245399 +21900,24.125,862.068966 +22200,24.453,920.245399 +22500,24.797,874.635569 +22800,25.150,854.700855 +23100,25.449,1010.101010 +23400,25.751,996.677741 +23700,26.066,952.380952 +24000,26.371,986.842105 +24300,26.703,903.614458 +24600,27.046,879.765396 +24900,27.360,958.466454 +25200,27.697,895.522388 +25500,28.033,917.431193 +25800,28.371,920.245399 +26100,28.729,840.336134 +26400,29.066,895.522388 +26700,29.384,946.372240 +27000,29.716,906.344411 +27300,30.070,852.272727 +27600,30.386,949.367089 +27900,30.707,940.438871 +28200,31.024,949.367089 +28500,31.332,977.198697 +28800,31.645,958.466454 +29100,31.991,869.565217 +29400,32.333,879.765396 +29700,32.674,884.955752 +30000,33.012,890.207715 +30300,33.354,882.352941 +30600,33.665,970.873786 +30900,33.995,911.854103 +31200,34.335,887.573964 +31500,34.671,895.522388 +31800,35.014,877.192982 +32100,35.351,892.857143 +32400,35.690,890.207715 +32700,36.038,862.068966 +33000,36.385,869.565217 +33300,36.711,928.792570 +33600,37.057,869.565217 +33900,37.391,900.900901 +34200,37.733,879.765396 +34500,38.067,900.900901 +34800,38.423,842.696629 +35100,38.750,923.076923 +35400,39.110,833.333333 +35700,39.443,906.344411 +36000,39.771,923.076923 +36300,40.105,900.900901 +36600,40.456,857.142857 +36900,40.782,925.925926 +37200,41.086,990.099010 +37500,41.390,990.099010 +37800,41.710,958.466454 +38100,42.017,980.392157 +38400,42.338,937.500000 +38700,42.650,961.538462 +39000,42.992,879.765396 +39300,43.332,887.573964 +39600,43.656,931.677019 +39900,43.991,895.522388 +40200,44.332,882.352941 +40500,44.641,970.873786 +40800,44.992,857.142857 +41100,45.321,917.431193 +41400,45.662,879.765396 diff --git a/outputData/NoGrid_MedDensity_128.csv b/outputData/NoGrid_MedDensity_128.csv new file mode 100644 index 0000000..0e555d6 --- /dev/null +++ b/outputData/NoGrid_MedDensity_128.csv @@ -0,0 +1,61 @@ +300,0.260,1153.846154 +600,0.416,1935.483871 +900,1.003,511.073254 +1200,1.755,400.000000 +1500,2.522,392.156863 +1800,3.303,384.615385 +2100,4.064,395.256917 +2400,4.848,382.653061 +2700,5.611,394.736842 +3000,6.392,385.109114 +3300,7.156,393.700787 +3600,7.935,391.644909 +3900,8.714,389.105058 +4200,9.491,386.597938 +4500,10.272,385.109114 +4800,11.039,391.644909 +5100,11.808,390.625000 +5400,12.582,388.601036 +5700,13.348,392.156863 +6000,14.121,394.736842 +6300,14.900,386.100386 +6600,15.684,384.615385 +6900,16.450,392.156863 +7200,17.219,390.625000 +7500,18.004,383.141762 +7800,18.784,385.109114 +8100,19.556,389.610390 +8400,20.337,385.109114 +8700,21.116,386.100386 +9000,21.906,386.597938 +9300,22.691,383.141762 +9600,23.462,390.117035 +9900,24.249,382.165605 +10200,25.026,387.596899 +10500,25.807,384.615385 +10800,26.574,392.156863 +11100,27.353,385.604113 +11400,28.131,386.100386 +11700,28.909,386.597938 +12000,29.664,398.406375 +12300,30.431,391.644909 +12600,31.174,404.858300 +12900,31.946,388.601036 +13200,32.713,392.670157 +13500,33.496,384.122919 +13800,34.272,388.098318 +14100,35.059,382.653061 +14400,35.838,385.604113 +14700,36.619,385.109114 +15000,37.389,390.625000 +15300,38.166,392.670157 +15600,38.942,389.610390 +15900,39.728,383.141762 +16200,40.499,390.117035 +16500,41.266,391.644909 +16800,42.039,389.610390 +17100,42.838,384.615385 +17400,43.603,392.670157 +17700,44.386,383.141762 +18000,45.172,382.653061 +18300,45.943,389.610390 diff --git a/outputData/NoGrid_MedDensity_32.csv b/outputData/NoGrid_MedDensity_32.csv new file mode 100644 index 0000000..8491720 --- /dev/null +++ b/outputData/NoGrid_MedDensity_32.csv @@ -0,0 +1,57 @@ +300,0.175,1714.285714 +600,0.324,2013.422819 +900,0.977,460.829493 +1200,1.740,394.736842 +1500,2.501,394.736842 +1800,3.270,390.625000 +2100,4.032,395.256917 +2400,4.794,395.256917 +2700,5.560,392.670157 +3000,6.327,391.644909 +3300,7.093,392.670157 +3600,7.862,391.134289 +3900,8.624,394.218134 +4200,9.387,393.700787 +4500,10.145,396.825397 +4800,10.914,395.256917 +5100,11.690,388.601036 +5400,12.465,387.596899 +5700,13.232,392.156863 +6000,13.994,394.218134 +6300,14.780,382.653061 +6600,15.545,393.184797 +6900,16.313,391.134289 +7200,17.067,398.406375 +7500,17.836,390.625000 +7800,18.614,386.597938 +8100,19.388,388.098318 +8400,20.170,384.615385 +8700,20.927,396.825397 +9000,21.709,392.670157 +9300,22.481,389.105058 +9600,23.249,391.644909 +9900,24.030,385.109114 +10200,24.807,387.096774 +10500,25.569,394.218134 +10800,26.343,393.184797 +11100,27.123,386.100386 +11400,27.886,394.218134 +11700,28.655,391.134289 +12000,29.414,397.350993 +12300,30.177,393.184797 +12600,30.932,398.406375 +12900,31.685,398.936170 +13200,32.446,395.256917 +13500,33.223,386.597938 +13800,33.993,390.625000 +14100,34.762,390.625000 +14400,35.538,387.596899 +14700,36.313,388.098318 +15000,37.076,394.218134 +15300,37.848,390.117035 +15600,38.622,388.601036 +15900,39.392,390.117035 +16200,40.162,390.625000 +16500,40.935,388.098318 +16800,41.698,394.736842 +17100,42.471,388.601036 diff --git a/outputData/NoGrid_MedDensity_512.csv b/outputData/NoGrid_MedDensity_512.csv new file mode 100644 index 0000000..4b84f57 --- /dev/null +++ b/outputData/NoGrid_MedDensity_512.csv @@ -0,0 +1,58 @@ +300,0.172,1744.186047 +600,0.301,2325.581395 +900,1.064,393.184797 +1200,1.883,367.197062 +1500,2.698,368.550369 +1800,3.513,370.370370 +2100,4.318,373.599004 +2400,5.132,369.003690 +2700,5.948,368.098160 +3000,6.765,368.098160 +3300,7.583,367.647059 +3600,8.404,366.300366 +3900,9.229,364.077670 +4200,10.045,368.550369 +4500,10.872,363.636364 +4800,11.679,372.670807 +5100,12.503,364.963504 +5400,13.315,370.370370 +5700,14.126,370.828183 +6000,14.948,368.550369 +6300,15.768,371.747212 +6600,16.573,374.064838 +6900,17.364,379.266751 +7200,18.176,369.913687 +7500,18.994,367.647059 +7800,19.814,366.300366 +8100,20.641,363.636364 +8400,21.446,373.599004 +8700,22.275,362.756953 +9000,23.093,367.647059 +9300,23.905,369.913687 +9600,24.721,368.098160 +9900,25.537,369.003690 +10200,26.363,366.748166 +10500,27.174,370.828183 +10800,27.991,372.670807 +11100,28.810,367.647059 +11400,29.602,379.266751 +11700,30.418,373.599004 +12000,31.219,375.939850 +12300,32.025,372.670807 +12600,32.832,372.670807 +12900,33.655,365.408039 +13200,34.480,364.077670 +13500,35.295,369.003690 +13800,36.112,368.098160 +14100,36.932,366.300366 +14400,37.743,370.370370 +14700,38.555,370.370370 +15000,39.387,361.445783 +15300,40.203,368.550369 +15600,41.007,374.064838 +15900,41.824,368.098160 +16200,42.650,364.077670 +16500,43.466,368.550369 +16800,44.279,369.913687 +17100,45.094,369.003690 +17400,45.905,370.370370 diff --git a/outputData/UniformGrid_HighDensity_128.csv b/outputData/UniformGrid_HighDensity_128.csv new file mode 100644 index 0000000..94ff202 --- /dev/null +++ b/outputData/UniformGrid_HighDensity_128.csv @@ -0,0 +1,62 @@ +300,0.698,429.799427 +600,1.337,469.483568 +900,1.996,455.927052 +1200,2.663,450.450450 +1500,3.356,433.526012 +1800,4.045,436.046512 +2100,4.739,432.276657 +2400,5.453,420.757363 +2700,6.154,429.184549 +3000,6.881,413.223140 +3300,7.595,420.168067 +3600,8.340,403.768506 +3900,9.071,410.958904 +4200,9.804,409.276944 +4500,10.541,407.608696 +4800,11.294,398.936170 +5100,12.042,402.144772 +5400,12.779,407.608696 +5700,13.521,404.312668 +6000,14.242,416.666667 +6300,14.976,409.276944 +6600,15.701,414.364641 +6900,16.436,408.719346 +7200,17.163,413.793103 +7500,17.880,418.994413 +7800,18.619,407.055631 +8100,19.359,405.953992 +8400,20.091,410.396717 +8700,20.824,409.836066 +9000,21.571,402.144772 +9300,22.316,403.225806 +9600,23.058,404.858300 +9900,23.807,401.069519 +10200,24.565,396.301189 +10500,25.324,395.256917 +10800,26.083,395.778364 +11100,26.844,395.256917 +11400,27.609,392.156863 +11700,28.362,398.936170 +12000,29.125,393.700787 +12300,29.887,394.218134 +12600,30.662,387.596899 +12900,31.416,398.406375 +13200,32.184,391.134289 +13500,32.955,389.610390 +13800,33.733,386.100386 +14100,34.506,388.601036 +14400,35.286,385.604113 +14700,36.060,388.098318 +15000,36.815,397.877984 +15300,37.573,396.301189 +15600,38.337,393.184797 +15900,39.102,392.670157 +16200,39.877,387.096774 +16500,40.653,387.596899 +16800,41.392,405.953992 +17100,42.165,388.601036 +17400,42.949,383.141762 +17700,43.696,402.144772 +18000,44.474,386.100386 +18300,45.246,389.105058 +18600,46.001,397.877984 diff --git a/outputData/UniformGrid_HighDensity_512.csv b/outputData/UniformGrid_HighDensity_512.csv new file mode 100644 index 0000000..14e75ba --- /dev/null +++ b/outputData/UniformGrid_HighDensity_512.csv @@ -0,0 +1,57 @@ +300,0.717,418.410042 +600,1.400,439.882698 +900,2.101,427.960057 +1200,2.810,423.728814 +1500,3.550,405.953992 +1800,4.290,405.953992 +2100,5.053,393.700787 +2400,5.822,390.625000 +2700,6.594,389.105058 +3000,7.356,394.218134 +3300,8.117,394.218134 +3600,8.893,387.596899 +3900,9.684,379.746835 +4200,10.460,387.096774 +4500,11.238,386.100386 +4800,12.036,376.411543 +5100,12.839,373.599004 +5400,13.657,367.197062 +5700,14.467,370.828183 +6000,15.270,374.064838 +6300,16.059,380.710660 +6600,16.838,385.604113 +6900,17.628,380.228137 +7200,18.411,383.631714 +7500,19.193,384.615385 +7800,20.010,367.647059 +8100,20.814,373.599004 +8400,21.590,387.096774 +8700,22.394,373.599004 +9000,23.188,378.310214 +9300,24.010,365.408039 +9600,24.837,362.756953 +9900,25.671,360.144058 +10200,26.506,359.712230 +10500,27.333,363.196126 +10800,28.144,370.370370 +11100,28.973,362.318841 +11400,29.795,365.408039 +11700,30.624,362.756953 +12000,31.447,364.963504 +12300,32.270,364.963504 +12600,33.125,351.288056 +12900,33.968,356.718193 +13200,34.801,360.576923 +13500,35.647,355.029586 +13800,36.494,355.029586 +14100,37.335,357.142857 +14400,38.174,358.422939 +14700,39.001,363.196126 +15000,39.829,362.756953 +15300,40.671,357.142857 +15600,41.510,357.568534 +15900,42.381,344.827586 +16200,43.229,354.191263 +16500,44.074,355.871886 +16800,44.935,348.837209 +17100,45.802,346.420323 diff --git a/outputData/UniformGrid_LowDensity_128.csv b/outputData/UniformGrid_LowDensity_128.csv new file mode 100644 index 0000000..570f16a --- /dev/null +++ b/outputData/UniformGrid_LowDensity_128.csv @@ -0,0 +1,147 @@ +300,0.352,852.272727 +600,0.660,977.198697 +900,0.966,983.606557 +1200,1.256,1038.062284 +1500,1.551,1020.408163 +1800,1.844,1027.397260 +2100,2.145,996.677741 +2400,2.435,1038.062284 +2700,2.725,1034.482759 +3000,3.017,1030.927835 +3300,3.314,1013.513514 +3600,3.612,1010.101010 +3900,3.906,1023.890785 +4200,4.205,1006.711409 +4500,4.505,1003.344482 +4800,4.806,996.677741 +5100,5.103,1013.513514 +5400,5.399,1013.513514 +5700,5.698,1006.711409 +6000,5.997,1006.711409 +6300,6.297,1000.000000 +6600,6.598,1000.000000 +6900,6.892,1023.890785 +7200,7.188,1016.949153 +7500,7.491,993.377483 +7800,7.792,1000.000000 +8100,8.095,993.377483 +8400,8.397,996.677741 +8700,8.700,990.099010 +9000,9.002,996.677741 +9300,9.305,993.377483 +9600,9.614,977.198697 +9900,9.916,996.677741 +10200,10.221,986.842105 +10500,10.522,996.677741 +10800,10.831,974.025974 +11100,11.138,980.392157 +11400,11.438,1003.344482 +11700,11.740,993.377483 +12000,12.039,1006.711409 +12300,12.350,967.741935 +12600,12.654,986.842105 +12900,12.959,990.099010 +13200,13.268,974.025974 +13500,13.573,986.842105 +13800,13.890,946.372240 +14100,14.194,990.099010 +14400,14.508,958.466454 +14700,14.819,967.741935 +15000,15.136,949.367089 +15300,15.438,996.677741 +15600,15.750,961.538462 +15900,16.058,977.198697 +16200,16.374,949.367089 +16500,16.693,943.396226 +16800,17.009,952.380952 +17100,17.315,980.392157 +17400,17.625,974.025974 +17700,17.932,977.198697 +18000,18.243,967.741935 +18300,18.552,974.025974 +18600,18.863,970.873786 +18900,19.167,986.842105 +19200,19.489,934.579439 +19500,19.798,970.873786 +19800,20.109,967.741935 +20100,20.411,993.377483 +20400,20.720,974.025974 +20700,21.036,949.367089 +21000,21.350,958.466454 +21300,21.657,980.392157 +21600,21.964,977.198697 +21900,22.278,961.538462 +22200,22.598,940.438871 +22500,22.908,970.873786 +22800,23.222,955.414013 +23100,23.532,974.025974 +23400,23.847,952.380952 +23700,24.150,993.377483 +24000,24.460,970.873786 +24300,24.769,977.198697 +24600,25.080,967.741935 +24900,25.392,964.630225 +25200,25.702,967.741935 +25500,26.018,952.380952 +25800,26.325,980.392157 +26100,26.643,946.372240 +26400,26.953,970.873786 +26700,27.257,990.099010 +27000,27.570,961.538462 +27300,27.884,958.466454 +27600,28.202,943.396226 +27900,28.519,949.367089 +28200,28.832,961.538462 +28500,29.146,955.414013 +28800,29.466,940.438871 +29100,29.776,970.873786 +29400,30.087,964.630225 +29700,30.397,970.873786 +30000,30.710,961.538462 +30300,31.040,909.090909 +30600,31.345,986.842105 +30900,31.684,887.573964 +31200,32.006,934.579439 +31500,32.315,974.025974 +31800,32.627,964.630225 +32100,32.938,967.741935 +32400,33.248,970.873786 +32700,33.561,961.538462 +33000,33.876,955.414013 +33300,34.196,940.438871 +33600,34.508,961.538462 +33900,34.817,974.025974 +34200,35.131,958.466454 +34500,35.446,955.414013 +34800,35.757,964.630225 +35100,36.069,964.630225 +35400,36.384,955.414013 +35700,36.699,952.380952 +36000,37.012,961.538462 +36300,37.326,961.538462 +36600,37.635,974.025974 +36900,37.948,961.538462 +37200,38.259,967.741935 +37500,38.574,955.414013 +37800,38.883,974.025974 +38100,39.196,961.538462 +38400,39.514,949.367089 +38700,39.834,940.438871 +39000,40.147,961.538462 +39300,40.455,977.198697 +39600,40.764,974.025974 +39900,41.077,961.538462 +40200,41.389,964.630225 +40500,41.705,952.380952 +40800,42.014,970.873786 +41100,42.329,955.414013 +41400,42.641,964.630225 +41700,42.952,964.630225 +42000,43.263,970.873786 +42300,43.580,946.372240 +42600,43.893,958.466454 +42900,44.212,943.396226 +43200,44.520,977.198697 +43500,44.827,983.606557 +43800,45.133,983.606557 +44100,45.450,949.367089 diff --git a/outputData/UniformGrid_MedDensity_128.csv b/outputData/UniformGrid_MedDensity_128.csv new file mode 100644 index 0000000..6416cc8 --- /dev/null +++ b/outputData/UniformGrid_MedDensity_128.csv @@ -0,0 +1,94 @@ +300,0.481,623.700624 +600,0.914,694.444444 +900,1.344,699.300699 +1200,1.773,700.934579 +1500,2.206,694.444444 +1800,2.637,697.674419 +2100,3.072,691.244240 +2400,3.514,680.272109 +2700,3.956,680.272109 +3000,4.400,677.200903 +3300,4.848,671.140940 +3600,5.296,669.642857 +3900,5.742,674.157303 +4200,6.193,666.666667 +4500,6.658,645.161290 +4800,7.126,643.776824 +5100,7.588,650.759219 +5400,8.049,650.759219 +5700,8.503,663.716814 +6000,8.954,665.188470 +6300,9.412,656.455142 +6600,9.875,649.350649 +6900,10.342,643.776824 +7200,10.810,642.398287 +7500,11.284,634.249471 +7800,11.757,635.593220 +8100,12.227,639.658849 +8400,12.690,649.350649 +8700,13.164,634.249471 +9000,13.637,635.593220 +9300,14.111,635.593220 +9600,14.589,628.930818 +9900,15.068,627.615063 +10200,15.530,652.173913 +10500,15.994,647.948164 +10800,16.461,642.398287 +11100,16.939,628.930818 +11400,17.424,619.834711 +11700,17.889,646.551724 +12000,18.370,625.000000 +12300,18.846,630.252101 +12600,19.323,630.252101 +12900,19.807,621.118012 +13200,20.283,630.252101 +13500,20.770,617.283951 +13800,21.232,650.759219 +14100,21.719,617.283951 +14400,22.205,618.556701 +14700,22.679,634.249471 +15000,23.155,630.252101 +15300,23.632,631.578947 +15600,24.124,609.756098 +15900,24.607,622.406639 +16200,25.088,626.304802 +16500,25.570,623.700624 +16800,26.061,612.244898 +17100,26.546,619.834711 +17400,27.027,625.000000 +17700,27.513,618.556701 +18000,27.987,634.249471 +18300,28.465,628.930818 +18600,28.946,625.000000 +18900,29.447,600.000000 +19200,29.915,642.398287 +19500,30.412,603.621730 +19800,30.896,621.118012 +20100,31.377,625.000000 +20400,31.876,602.409639 +20700,32.352,632.911392 +21000,32.837,618.556701 +21300,33.330,609.756098 +21600,33.814,619.834711 +21900,34.299,619.834711 +22200,34.778,626.304802 +22500,35.265,617.283951 +22800,35.757,610.997963 +23100,36.254,604.838710 +23400,36.740,618.556701 +23700,37.225,619.834711 +24000,37.704,627.615063 +24300,38.186,622.406639 +24600,38.660,634.249471 +24900,39.164,596.421471 +25200,39.662,603.621730 +25500,40.148,617.283951 +25800,40.647,603.621730 +26100,41.146,601.202405 +26400,41.644,604.838710 +26700,42.144,600.000000 +27000,42.653,591.715976 +27300,43.148,606.060606 +27600,43.660,587.084149 +27900,44.150,613.496933 +28200,44.660,589.390963 diff --git a/outputData/UniformGrid_MedDensity_32.csv b/outputData/UniformGrid_MedDensity_32.csv new file mode 100644 index 0000000..bfb1da5 --- /dev/null +++ b/outputData/UniformGrid_MedDensity_32.csv @@ -0,0 +1,72 @@ +300,0.640,468.750000 +600,1.208,529.100529 +900,1.786,519.930676 +1200,2.374,511.073254 +1500,2.985,491.803279 +1800,3.582,503.355705 +2100,4.176,505.902192 +2400,4.790,489.396411 +2700,5.394,497.512438 +3000,5.994,500.834725 +3300,6.600,495.867769 +3600,7.216,487.012987 +3900,7.826,492.610837 +4200,8.435,493.421053 +4500,9.045,492.610837 +4800,9.639,506.756757 +5100,10.254,488.599349 +5400,10.873,485.436893 +5700,11.491,486.223663 +6000,12.129,470.957614 +6300,12.755,480.000000 +6600,13.398,468.018721 +6900,14.019,483.870968 +7200,14.633,489.396411 +7500,15.256,481.540931 +7800,15.883,479.233227 +8100,16.523,469.483568 +8400,17.138,488.599349 +8700,17.772,473.933649 +9000,18.400,478.468900 +9300,19.030,476.947536 +9600,19.652,483.091787 +9900,20.265,489.396411 +10200,20.889,481.540931 +10500,21.524,473.933649 +10800,22.169,465.116279 +11100,22.787,487.012987 +11400,23.415,477.707006 +11700,24.046,476.190476 +12000,24.671,480.769231 +12300,25.301,476.947536 +12600,25.930,477.707006 +12900,26.578,463.678516 +13200,27.214,471.698113 +13500,27.849,473.186120 +13800,28.467,485.436893 +14100,29.105,470.957614 +14400,29.742,471.698113 +14700,30.389,464.396285 +15000,31.023,473.933649 +15300,31.653,476.947536 +15600,32.279,480.000000 +15900,32.919,469.483568 +16200,33.563,466.562986 +16500,34.197,473.933649 +16800,34.824,478.468900 +17100,35.453,477.707006 +17400,36.096,466.562986 +17700,36.743,464.396285 +18000,37.393,462.249615 +18300,38.028,473.186120 +18600,38.666,470.219436 +18900,39.318,461.538462 +19200,39.955,471.698113 +19500,40.589,473.933649 +19800,41.237,463.678516 +20100,41.886,462.962963 +20400,42.530,465.838509 +20700,43.156,480.769231 +21000,43.806,462.249615 +21300,44.443,470.957614 +21600,45.098,459.418070 diff --git a/outputData/UniformGrid_MedDensity_512.csv b/outputData/UniformGrid_MedDensity_512.csv new file mode 100644 index 0000000..0988ea9 --- /dev/null +++ b/outputData/UniformGrid_MedDensity_512.csv @@ -0,0 +1,91 @@ +300,0.509,589.390963 +600,0.966,656.455142 +900,1.421,660.792952 +1200,1.883,650.759219 +1500,2.344,650.759219 +1800,2.812,642.398287 +2100,3.288,631.578947 +2400,3.762,634.249471 +2700,4.242,626.304802 +3000,4.723,625.000000 +3300,5.206,622.406639 +3600,5.690,621.118012 +3900,6.180,613.496933 +4200,6.670,613.496933 +4500,7.170,600.000000 +4800,7.676,594.059406 +5100,8.171,607.287449 +5400,8.673,598.802395 +5700,9.169,606.060606 +6000,9.668,601.202405 +6300,10.169,600.000000 +6600,10.675,594.059406 +6900,11.180,595.238095 +7200,11.693,585.937500 +7500,12.200,592.885375 +7800,12.710,590.551181 +8100,13.212,598.802395 +8400,13.718,592.885375 +8700,14.222,597.609562 +9000,14.730,591.715976 +9300,15.242,587.084149 +9600,15.746,596.421471 +9900,16.253,592.885375 +10200,16.761,591.715976 +10500,17.272,588.235294 +10800,17.788,582.524272 +11100,18.298,589.390963 +11400,18.811,585.937500 +11700,19.324,585.937500 +12000,19.838,584.795322 +12300,20.362,573.613767 +12600,20.874,587.084149 +12900,21.385,587.084149 +13200,21.895,589.390963 +13500,22.413,580.270793 +13800,22.962,546.448087 +14100,23.471,591.715976 +14400,23.988,580.270793 +14700,24.506,580.270793 +15000,25.031,572.519084 +15300,25.545,585.937500 +15600,26.057,587.084149 +15900,26.570,585.937500 +16200,27.079,589.390963 +16500,27.599,578.034682 +16800,28.115,582.524272 +17100,28.634,579.150579 +17400,29.154,578.034682 +17700,29.671,581.395349 +18000,30.188,581.395349 +18300,30.707,579.150579 +18600,31.231,573.613767 +18900,31.756,573.613767 +19200,32.275,579.150579 +19500,32.797,575.815739 +19800,33.319,574.712644 +20100,33.847,569.259962 +20400,34.363,582.524272 +20700,34.888,572.519084 +21000,35.410,575.815739 +21300,35.940,568.181818 +21600,36.465,571.428571 +21900,36.994,568.181818 +22200,37.520,571.428571 +22500,38.041,576.923077 +22800,38.566,572.519084 +23100,39.089,574.712644 +23400,39.617,569.259962 +23700,40.136,579.150579 +24000,40.660,573.613767 +24300,41.176,582.524272 +24600,41.705,568.181818 +24900,42.237,564.971751 +25200,42.770,563.909774 +25500,43.302,564.971751 +25800,43.834,564.971751 +26100,44.370,560.747664 +26400,44.916,550.458716 +26700,45.458,554.528651 +27000,45.998,556.586271 +27300,46.529,566.037736 diff --git a/outputData/csvReader.py b/outputData/csvReader.py new file mode 100644 index 0000000..c35e884 --- /dev/null +++ b/outputData/csvReader.py @@ -0,0 +1,80 @@ +import sys +import os +import re +import numpy as np +import matplotlib.pyplot as plt +import csv + + +def readCSV(filename): + """ + takes in a filename + returns a numpy array, with the first row as the + timestamp in seconds, and the second row + as the fps across the last time block + """ + results = [] + with open(filename) as csv_file: + reader = csv.reader(csv_file, delimiter=',') + for line in reader: + timestamp = float(line[1]) + fps = float(line[2]) + results.append([timestamp, fps]) + + return np.array(results).T + +def makeGraphs(resultSets, title): + """ + Displays the resultant data sets, along with a given title + """ + fig, ax = plt.subplots(1) + for filename, data in resultSets: + ax.plot(data[0], data[1], label = cleanFileName(filename)) + + ax.legend() + plt.xlabel("Time (seconds)") + plt.ylabel("Ticks/Frames per Second") + + fig.suptitle(title) + fig.set_size_inches(10,6) + + #plt.show() #uncomment this to display the graph on your screen + filePath = makeSavePath(title) + plt.savefig(filePath) + +def cleanFileName(fileName): + """ + Turns the filename into a string appropriate for labeling a data series + """ + regstr = "([A-Z][a-z]+)([A-Z][a-z]+)_([A-Z][a-z]+)([A-Z][a-z]+)_([0-9]+).csv" + tokens = re.match(regstr, fileName).groups() + + retval = "%s %s, %s %s, BlockSize %s" % tokens + return retval + +def makeSavePath(title): + filePath = os.path.join("..", "images", title + ".png") + return filePath + + + +def main(): + if len(sys.argv) < 3: + print("Please input a title and file names") + exit(0) + + resultSets = [] + + for i, fileName in enumerate(sys.argv): + if i == 0 or i == 1: + continue + resultSets.append((fileName, readCSV(fileName))) + + makeGraphs(resultSets, sys.argv[1]) + + + +if __name__ == "__main__": + main() + + diff --git a/src/kernel.cu b/src/kernel.cu index 74dffcb..8fb1b87 100644 --- a/src/kernel.cu +++ b/src/kernel.cu @@ -3,9 +3,15 @@ #include #include #include +#include +#include +#include #include "utilityCore.hpp" +#include "cudaMat4.hpp" #include "kernel.h" +#define DEBUGOUT 0 + // LOOK-2.1 potentially useful for doing grid-based neighbor search #ifndef imax #define imax( a, b ) ( ((a) > (b)) ? (a) : (b) ) @@ -76,9 +82,9 @@ glm::vec3 *dev_vel2; // For efficient sorting and the uniform grid. These should always be parallel. int *dev_particleArrayIndices; // What index in dev_pos and dev_velX represents this particle? int *dev_particleGridIndices; // What grid cell is this particle in? -// needed for use with thrust -thrust::device_ptr dev_thrust_particleArrayIndices; -thrust::device_ptr dev_thrust_particleGridIndices; + +cudaMat3* dev_pvv;//for storing the pos/vel1/vel2 all in the same place + int *dev_gridCellStartIndices; // What part of dev_particleArrayIndices belongs int *dev_gridCellEndIndices; // to this cell? @@ -92,8 +98,50 @@ int gridCellCount; int gridSideCount; float gridCellWidth; float gridInverseCellWidth; + +__constant__ int dgridCellCount = -1; +__constant__ int dgridSideCount = -1; +__constant__ float dgridCellWidth = -1.0f; +__constant__ float dgridInverseCellWidth = -1.0f; + glm::vec3 gridMinimum; +/******************************** +* Forward Function declarations +*********************************/ + +/** +* Given a location in our world, gives us back which grid cell contains that point +* Additionally, takes in an optional pointer to a spot into which to throw +* the octant of our point within that grid cube +* If passed NULL, the function will not determine the octant. +*/ +__device__ int locToGridIndex(glm::vec3* pos, int* octant); + +/** +* Converts a linear grid index to a trio of grid coordinates +*/ +__device__ glm::ivec3 gridIndexToGridCoords(int gridIndex); +/** +* Converts a trio of grid coordinates to a linear grid index +*/ +__device__ int gridCoordsToGridIndex(glm::ivec3 gridCoords); +/** +* These fill in an array `results` with the indexes of cells adjacent to the provided one +* This list does include the provided cell as well +* Notably, this requires the octant of the "center" as well +* This NEEDS to be passed an int array with space for 8 result numbers +*/ +__device__ void getNeighboringCellIndexes(int gridIndex, int octant, int* results); +__device__ void getNeighboringCellIndexesTrio(glm::ivec3 gridCoords, int octant, int* results); + +/** +* Function for finding the beginning/end of cellNumber in the sorted list gridIndices +* Returns a 2d vector with the (indexBeginning, indexEnd) values (the latter exclusive), or (-1, -1) if not found +* So, given a `cellNumber`, the array `gridIndices` contains that number from `indexBeginning` through `indexEnd -1` +*/ +__device__ glm::ivec2 findMyGridIndices(int cellNumber, int* gridIndices, int N); + /****************** * initSimulation * ******************/ @@ -168,7 +216,28 @@ void Boids::initSimulation(int N) { gridMinimum.y -= halfGridWidth; gridMinimum.z -= halfGridWidth; + //Copy some of these grid constants to device memory, to abuse later + cudaMemcpyToSymbol(dgridCellCount, &gridCellCount, sizeof(int)); + cudaMemcpyToSymbol(dgridSideCount, &gridSideCount, sizeof(int)); + cudaMemcpyToSymbol(dgridCellWidth, &gridCellWidth, sizeof(float)); + cudaMemcpyToSymbol(dgridInverseCellWidth, &gridInverseCellWidth, sizeof(float)); + + printf("halfSideCount %d, gridSideCount %d, gridCellCount %d, gridInverseCellWidth %0.1f, halfGridWidth %0.1f, gridmin (x, y, z) (%.1f, %.1f, %.1f)\n", + halfSideCount, gridSideCount, gridCellCount, gridInverseCellWidth, halfGridWidth, gridMinimum.x, gridMinimum.y, gridMinimum.z); + // TODO-2.1 TODO-2.3 - Allocate additional buffers here. + + 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_particleGridIndices, N * sizeof(int)); + checkCUDAErrorWithLine("cudaMalloc dev_particleGridIndices failed!"); + cudaMalloc((void**)& dev_particleArrayIndices, N * sizeof(int)); + checkCUDAErrorWithLine("cudaMalloc dev_particleArrayIndices failed!"); + + cudaMalloc((void**)& dev_pvv, N * sizeof(cudaMat3)); + cudaDeviceSynchronize(); } @@ -210,19 +279,305 @@ __global__ void kernCopyVelocitiesToVBO(int N, glm::vec3 *vel, float *vbo, float void Boids::copyBoidsToVBO(float *vbodptr_positions, float *vbodptr_velocities) { dim3 fullBlocksPerGrid((numObjects + blockSize - 1) / blockSize); - kernCopyPositionsToVBO << > >(numObjects, dev_pos, vbodptr_positions, scene_scale); - kernCopyVelocitiesToVBO << > >(numObjects, dev_vel1, vbodptr_velocities, scene_scale); - - checkCUDAErrorWithLine("copyBoidsToVBO failed!"); + kernCopyPositionsToVBO <<< fullBlocksPerGrid, blockSize >>>(numObjects, dev_pos, vbodptr_positions, scene_scale); + checkCUDAErrorWithLine("copyBoidsToVBO failed on Pos!"); + kernCopyVelocitiesToVBO <<< fullBlocksPerGrid, blockSize >>>(numObjects, dev_vel1, vbodptr_velocities, scene_scale); + checkCUDAErrorWithLine("copyBoidsToVBO failed on Vel!"); cudaDeviceSynchronize(); } +/****************** +* math helpers * +******************/ + +/** +* Helper function to compute distance between two boids +* Added as the first steps of "can I modify this program sensibly" +*/ +__device__ float computeDistance(const glm::vec3* pos1, const glm::vec3* pos2){ + double result = sqrt((pos2->x - pos1->x) * (pos2->x - pos1->x) + + (pos2->y - pos1->y) * (pos2->y - pos1->y) + + (pos2->z - pos1->z) * (pos2->z - pos1->z)); + //double result = glm::distance(*pos1, *pos2); + return result; + +}//kernComputeDistance + +/** +* Clamps the speed down to maxSpeed +* +* Does so in-place (DOES modify data) +*/ +__device__ void clampSpeed(glm::vec3* vel){ + glm::vec3 zeroPoint = glm::vec3(0.0f, 0.0f, 0.0f); + double curSpeed = computeDistance(&zeroPoint, vel); + if (curSpeed > maxSpeed){ + double scaleFactor = maxSpeed / curSpeed; + *vel *= scaleFactor; + }//if + +}//clampSpeed + +/******************** +* grid math helpers * +*********************/ + +__device__ int locToGridIndex(glm::vec3* pos, int* octant) { + //positions range between -scene_scale and scene_scale + float gridInverseCellWidth = dgridInverseCellWidth; + float xGridPos = pos->x * gridInverseCellWidth;//range between, for our example, -10 and 10 + float yGridPos = pos->y * gridInverseCellWidth; + float zGridPos = pos->z * gridInverseCellWidth; + + glm::ivec3 gridTrio = glm::ivec3((int)(glm::floor(xGridPos)), (int)(glm::floor(yGridPos)), (int)(glm::floor(zGridPos))); + int retval = gridCoordsToGridIndex(gridTrio); + + if (octant != NULL) { + uint8_t topx = (((int)xGridPos) == ((int)(xGridPos - 0.5))); + uint8_t topy = (((int)yGridPos) == ((int)(yGridPos - 0.5))); + uint8_t topz = (((int)zGridPos) == ((int)(zGridPos - 0.5))); + + *octant = (int)(topz << 2 | topy < 1 | topx); + }//if + return retval; +}//locToGridIndex + +__device__ glm::ivec3 gridIndexToGridCoords(int gridIndex) { + int gridSideCount = dgridSideCount; + int coordz = gridIndex / (gridSideCount * gridSideCount); + int coordy = (gridIndex / gridSideCount) % gridSideCount; + int coordx = gridIndex % gridSideCount; + + glm::ivec3 retval = glm::ivec3(coordx, coordy, coordz); + retval -= (gridSideCount / 2); + + return retval; + +}//gridIndexToGridCoords + +// 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? +// for(x) +// for(y) +// for(z)? Or some other order? +__device__ int gridCoordsToGridIndex(glm::ivec3 gridCoords) { + int gridSideCount = dgridSideCount; + int halfSideCount = gridSideCount / 2; + int coordx = gridCoords.x + halfSideCount; + int coordy = gridCoords.y + halfSideCount; + int coordz = gridCoords.z + halfSideCount; + int retval = coordx + (coordy * gridSideCount) + (coordz * gridSideCount * gridSideCount); + return retval; +}//gridCoordsToGridIndex + +__device__ void getNeighboringCellIndexes(int gridIndex, int octant, int* results) { + glm::ivec3 gridCoord = gridIndexToGridCoords(gridIndex); + getNeighboringCellIndexesTrio(gridCoord, octant, results); +}//getNeighboringCellIndexes + +__device__ void getNeighboringCellIndexesTrio(glm::ivec3 gridCoords, int octant, int* results) { + int myIndex = gridCoordsToGridIndex(gridCoords); + results[0] = myIndex;//check own cell for sure + + //stand-in for booleans because I don't trust the conversions + uint8_t xtop = (uint8_t)(octant & 0x01); + uint8_t ytop = (uint8_t)(octant & 0x02); + uint8_t ztop = (uint8_t)(octant & 0x04); + int x = gridCoords.x; + int y = gridCoords.y; + int z = gridCoords.z; + + results[1] = gridCoordsToGridIndex(glm::ivec3(xtop ? x + 1 : x - 1, + y, + z)); + results[2] = gridCoordsToGridIndex(glm::ivec3(x, + ytop ? y + 1 : y - 1, + z)); + results[3] = gridCoordsToGridIndex(glm::ivec3(x, + y, + ztop ? z + 1 : z - 1)); + results[4] = gridCoordsToGridIndex(glm::ivec3(xtop ? x + 1 : x - 1, + ytop ? y + 1 : y - 1, + z)); + results[5] = gridCoordsToGridIndex(glm::ivec3(xtop ? x + 1 : x - 1, + y, + ztop ? z + 1 : z - 1)); + results[6] = gridCoordsToGridIndex(glm::ivec3(x, + ytop ? y + 1 : y - 1, + ztop ? z + 1 : z - 1)); + results[7] = gridCoordsToGridIndex(glm::ivec3(xtop ? x + 1 : x - 1, + ytop ? y + 1 : y - 1, + ztop ? z + 1 : z - 1)); + +}//getNeighboringCellIndexesTrio + +__device__ glm::ivec2 findMyGridIndices(int cellNumber, int* gridIndices, int N){ + int i = 0; + int beginIndex = -1; + int endIndex = -1; + while (i < N) { + if (beginIndex == -1 && gridIndices[i] == cellNumber) { + beginIndex = i; + }//if + else if (beginIndex != -1 && endIndex == -1 && gridIndices[i] != cellNumber) { + endIndex = i; + break;//for what good it does us, efficiency-wise, to break the loop early in the land of warps + }//else + else if (beginIndex != -1 && endIndex == -1 && i == N) { + endIndex = N; + }//I don't think it ever hits here? + i++; + }//while + if (beginIndex > endIndex) { + endIndex = N;//make sure we don't run off the end + }//if + if (beginIndex > endIndex) { + printf("How is %d still more than %d?\n", beginIndex, endIndex); + }//if + + return glm::ivec2(beginIndex, endIndex); + +}//findMyGridIndices + +/****************** +* mem xfers * +******************/ + +/** +* Puts the pos, vel1, and vel2 data into the dev_pvv array +*/ +__global__ void pvvConverge(int N, cudaMat3* pvv, glm::vec3* pos, glm::vec3* vel1, glm::vec3* vel2) { + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= N) { + return; + } + + pvv[index] = { glm::vec3(pos[index]), glm::vec3(vel1[index]), glm::vec3(vel2[index]) }; +}//pvvConverge + +/** +* Puts the pvv data into the dev_pos, dev_vel1, and dev_vel2 arrays +*/ +__global__ void pvvSplit(int N, cudaMat3* pvv, glm::vec3* pos, glm::vec3* vel1, glm::vec3* vel2) { + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= N) { + return; + } + glm::vec3 myPos = glm::vec3(pvv[index].x); + glm::vec3 myVel1 = glm::vec3(pvv[index].y); + glm::vec3 myVel2 = glm::vec3(pvv[index].z); + + pos[index] = glm::vec3(pvv[index].x); + vel1[index] = glm::vec3(pvv[index].y); + vel2[index] = glm::vec3(pvv[index].z); +}//pvvSplit /****************** * stepSimulation * ******************/ +/** +* Given a self-boid and another single boid, returns the velocity contribution of the pair for rule 2 +* Will call this method for multiple boids +*/ +__device__ glm::vec3 computeRule2VelContributionSingle(const glm::vec3* myPos, const glm::vec3* theirPos, + const glm::vec3* myVel, const glm::vec3* theirVel){ + float distBetween = computeDistance(myPos, theirPos); + if (distBetween > rule2Distance){ + return glm::vec3(0.0f, 0.0f, 0.0f); + }//if + + glm::vec3 distVec = *myPos - *theirPos; + + return distVec; + +}//computeRule2VelContributionSingle + +/** +* Given a self-boid and another single boid, returns the velocity contribution of the pair for rule 3 +* Will call this method for multiple boids +*/ +__device__ glm::vec3 computeRule3VelContributionSingle(const glm::vec3* myPos, const glm::vec3* theirPos, + const glm::vec3* myVel, const glm::vec3* theirVel){ + float distBetween = computeDistance(myPos, theirPos); + if (distBetween > rule3Distance){ + return glm::vec3(0.0f, 0.0f, 0.0f); + }//if + + return *theirVel; + +}//computeRule3VelContributionSingle + +/** +* Calculates the velocity contribution from rule 1 for all boids +*/ +__device__ glm::vec3 computeRule1VelContribution(int N, int iSelf, const glm::vec3* pos, const glm::vec3* vel){ + + glm::vec3 perceivedCenter = glm::vec3(0.0f, 0.0f, 0.0f); + int numNeighbors = 0; + + glm::vec3 myPos = pos[iSelf]; + //glm::vec3 myVel = vel[iSelf]; + + for (int i = 0; i < N; i++){ + if (i == iSelf) continue; + if (computeDistance(&myPos, &pos[i]) < rule1Distance){ + numNeighbors++; + perceivedCenter += pos[i]; + }//if a neighbor + + }//for each boid + + if (numNeighbors < 1) return glm::vec3(0.0f, 0.0f, 0.0f); + + perceivedCenter /= numNeighbors; + glm::vec3 resultVector = perceivedCenter - myPos; + resultVector *= rule1Scale; + + return resultVector; + +}//computeRule1VelContribution + +/** +* Calculates the velocity contribution from rule 2 for all boids +*/ +__device__ glm::vec3 computeRule2VelContribution(int N, int iSelf, const glm::vec3* pos, const glm::vec3* vel){ + + glm::vec3 velChange = glm::vec3(0.0f, 0.0f, 0.0f); + + glm::vec3 myPos = pos[iSelf]; + glm::vec3 myVel = vel[iSelf]; + + for (int i = 0; i < N; i++){ + if (i == iSelf) continue; + velChange += computeRule2VelContributionSingle(&myPos, &pos[i], &myVel, &vel[i]); + }//for each boid + + return velChange * rule2Scale; + +}//computeRule2VelContribution + +/** +* Calculates the velocity contribution from rule 3 for all boids +*/ +__device__ glm::vec3 computeRule3VelContribution(int N, int iSelf, const glm::vec3* pos, const glm::vec3* vel){ + + glm::vec3 velChange = glm::vec3(0.0f, 0.0f, 0.0f); + + glm::vec3 myPos = pos[iSelf]; + glm::vec3 myVel = vel[iSelf]; + + for (int i = 0; i < N; i++){ + if (i == iSelf) continue; + velChange += computeRule3VelContributionSingle(&myPos, &pos[i], &myVel, &vel[i]); + }//for each boid + + return velChange * rule3Scale; + +}//computeRule3VelContribution + /** * LOOK-1.2 You can use this as a helper for kernUpdateVelocityBruteForce. * __device__ code can be called from a __global__ context @@ -230,21 +585,35 @@ 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); -} + // Rule 1: boids fly towards their local perceived center of mass, which excludes themselves + glm::vec3 rule1VelChange = computeRule1VelContribution(N, iSelf, pos, vel); + // Rule 2: boids try to stay a distance d away from each other + glm::vec3 rule2VelChange = computeRule2VelContribution(N, iSelf, pos, vel); + // Rule 3: boids try to match the speed of surrounding boids + glm::vec3 rule3VelChange = computeRule3VelContribution(N, iSelf, pos, vel); + + return rule1VelChange + rule2VelChange + rule3VelChange; + +}//computeVelocityChange /** * TODO-1.2 implement basic flocking * For each of the `N` bodies, update its position based on its current velocity. */ __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? + glm::vec3 *vel1, glm::vec3 *vel2) { + + int index = threadIdx.x + (blockIdx.x * blockDim.x); + + // Compute a new velocity based on pos and vel1 + glm::vec3 velChange = computeVelocityChange(N, index, pos, vel1); + + // Clamp the speed + glm::vec3 newVel = vel1[index] + velChange; + clampSpeed(&newVel); + + // Record the new velocity into vel2. Question: why NOT vel1? + vel2[index] = newVel; } /** @@ -272,25 +641,56 @@ __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? -// for(x) -// 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; -} - -__global__ void kernComputeIndices(int N, int gridResolution, - glm::vec3 gridMin, float inverseCellWidth, +__global__ void kernComputeIndices(int N, float inverseCellWidth, glm::vec3 *pos, int *indices, int *gridIndices) { - // TODO-2.1 - // - 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; + } + + //dev_particleGridIndices + gridIndices[index] = locToGridIndex(&pos[index], NULL); + //dev_particleArrayIndices + indices[index] = index; + } +void Boids::sortGridIndices(int N, bool coherentGrid) { + + + thrust::device_ptr dev_thrust_particleGridIndices(dev_particleGridIndices); + + if (coherentGrid) { + thrust::device_ptr dev_thrust_pvv(dev_pvv); + thrust::sort_by_key(dev_thrust_particleGridIndices, dev_thrust_particleGridIndices + N, dev_thrust_pvv); + }//if + else { + thrust::device_ptr dev_thrust_particleArrayIndices(dev_particleArrayIndices); + thrust::sort_by_key(dev_thrust_particleGridIndices, dev_thrust_particleGridIndices + N, dev_thrust_particleArrayIndices); + }//else + +}//sortGridIndices + +/** +* Fills in the dev_gridCellStartIndices and dev_gridCellEndIndices arrays +* This is happening once for each cell, not once for each boid +*/ +__global__ void kernFindGridStartEnds(int N, int* particleGridIndices, int* gridStarts, int* gridEnds) { + // TODO-2.1 + // 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 gridCellCount = dgridCellCount; + int index = (blockIdx.x * blockDim.x) + threadIdx.x; + if (index >= gridCellCount) return; + + glm::ivec2 startEnd = findMyGridIndices(index, particleGridIndices, N); + gridStarts[index] = startEnd.x; + gridEnds[index] = startEnd.y; + +}//kernFindGridStartEnds + // 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) { @@ -300,20 +700,80 @@ __global__ void kernResetIntBuffer(int N, int *intBuffer, int value) { } } -__global__ void kernIdentifyCellStartEnd(int N, int *particleGridIndices, - int *gridCellStartIndices, int *gridCellEndIndices) { - // TODO-2.1 - // 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!" -} - __global__ void kernUpdateVelNeighborSearchScattered( - int N, int gridResolution, glm::vec3 gridMin, - float inverseCellWidth, float cellWidth, - int *gridCellStartIndices, int *gridCellEndIndices, + int N, int *gridCellStartIndices, int *gridCellEndIndices, int *particleArrayIndices, glm::vec3 *pos, glm::vec3 *vel1, glm::vec3 *vel2) { + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= N) { + return; + }//if + + //find our grid cell and octant + glm::vec3 myPos = pos[index]; + glm::vec3 myVel = vel1[index]; + int octant = -1;//will get overwritten + int myGridIndex = locToGridIndex(&myPos, &octant); + + //get the grid indices for our neighbors + int neighborGridIndices[8] = { -1, -1, -1, -1, -1, -1, -1, -1 }; + getNeighboringCellIndexes(myGridIndex, octant, neighborGridIndices); + + //Accumulators for velocity contributions + glm::vec3 rule1VelChange = glm::vec3(0.0f, 0.0f, 0.0f); + glm::vec3 rule2VelChange = glm::vec3(0.0f, 0.0f, 0.0f); + glm::vec3 rule3VelChange = glm::vec3(0.0f, 0.0f, 0.0f); + glm::vec3 rule1Center = glm::vec3(0.0f, 0.0f, 0.0f); + int rule1NumNeighbors = 0; + + //for each cell that could be close to us + for (int i = 0; i < 8; i++) { + int neighborCellIndex = neighborGridIndices[i]; + if (neighborCellIndex < 0) continue;//ignoring a weird fringe error! + int cellStart = gridCellStartIndices[neighborCellIndex]; + int cellEnd = gridCellEndIndices[neighborCellIndex]; + //for each boid in that cell + int j = cellStart; + while (j != -1 && j != cellEnd && j < N) { + int neighborBoidIndex = particleArrayIndices[j]; + if (neighborBoidIndex != index) { + glm::vec3 theirPos = pos[neighborBoidIndex]; + glm::vec3 theirVel = vel1[neighborBoidIndex]; + //do the contribution math + //Rule 1 + if (computeDistance(&myPos, &theirPos) < rule1Distance) { + rule1NumNeighbors++; + rule1Center += theirPos; + }//if + //Rule 2 + rule2VelChange += computeRule2VelContributionSingle(&myPos, &theirPos, + &myVel, &theirVel); + //Rule 3 + rule3VelChange += computeRule3VelContributionSingle(&myPos, &theirPos, + &myVel, &theirVel); + //increment j, to find other boids in that cell + }//if + + j++; + }//while + }//for + if (rule1NumNeighbors > 0) { + rule1Center /= rule1NumNeighbors; + rule1VelChange = rule1Center - myPos; + }//if we have neighbors for rule 1 + + + //scale the velocity contributions + rule1VelChange *= rule1Scale; + rule2VelChange *= rule2Scale; + rule3VelChange *= rule3Scale; + + //Add up vel contributions + glm::vec3 newVel = myVel + rule1VelChange + rule2VelChange + rule3VelChange; + clampSpeed(&newVel); + + vel2[index] = newVel; + // 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 @@ -325,10 +785,77 @@ __global__ void kernUpdateVelNeighborSearchScattered( } __global__ void kernUpdateVelNeighborSearchCoherent( - int N, int gridResolution, glm::vec3 gridMin, - float inverseCellWidth, float cellWidth, + int N, int *gridCellStartIndices, int *gridCellEndIndices, - glm::vec3 *pos, glm::vec3 *vel1, glm::vec3 *vel2) { + cudaMat3* pvv) { + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= N) { + return; + }//if + + glm::vec3 myPos = pvv[index].x; + glm::vec3 myVel = pvv[index].y; + int octant = -1; + int myGridIndex = locToGridIndex(&myPos, &octant); + + //get the grid indices for our neighbors + int neighborGridIndices[8] = { -1, -1, -1, -1, -1, -1, -1, -1 }; + getNeighboringCellIndexes(myGridIndex, octant, neighborGridIndices); + + //Accumulators for velocity contributions + glm::vec3 rule1VelChange = glm::vec3(0.0f, 0.0f, 0.0f); + glm::vec3 rule2VelChange = glm::vec3(0.0f, 0.0f, 0.0f); + glm::vec3 rule3VelChange = glm::vec3(0.0f, 0.0f, 0.0f); + glm::vec3 rule1Center = glm::vec3(0.0f, 0.0f, 0.0f); + int rule1NumNeighbors = 0; + + //for each cell that could be close to us + for (int i = 0; i < 8; i++) { + int neighborCellIndex = neighborGridIndices[i]; + if (neighborCellIndex < 0) continue;//ignoring a weird fringe error! + int cellStart = gridCellStartIndices[neighborCellIndex]; + int cellEnd = gridCellEndIndices[neighborCellIndex]; + //for each boid in that cell + int j = cellStart; + while (j != -1 && j != cellEnd && j < N) { + int neighborBoidIndex = j;//this is a change from the scattered, and redundant, but may be good for debug + if (neighborBoidIndex != index) { + glm::vec3 theirPos = pvv[neighborBoidIndex].x; + glm::vec3 theirVel = pvv[neighborBoidIndex].y; + //do the contribution math + //Rule 1 + if (computeDistance(&myPos, &theirPos) < rule1Distance) { + rule1NumNeighbors++; + rule1Center += theirPos; + }//if + //Rule 2 + rule2VelChange += computeRule2VelContributionSingle(&myPos, &theirPos, + &myVel, &theirVel); + //Rule 3 + rule3VelChange += computeRule3VelContributionSingle(&myPos, &theirPos, + &myVel, &theirVel); + //increment j, to find other boids in that cell + }//if + + j++; + }//while + }//for + if (rule1NumNeighbors > 0) { + rule1Center /= rule1NumNeighbors; + rule1VelChange = rule1Center - myPos; + }//if we have neighbors for rule 1 + + //scale the velocity contributions + rule1VelChange *= rule1Scale; + rule2VelChange *= rule2Scale; + rule3VelChange *= rule3Scale; + + //Add up vel contributions + glm::vec3 newVel = myVel + rule1VelChange + rule2VelChange + rule3VelChange; + clampSpeed(&newVel); + + pvv[index].z = newVel; + // 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 @@ -347,23 +874,51 @@ __global__ void kernUpdateVelNeighborSearchCoherent( * Step the entire N-body simulation by `dt` seconds. */ 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 + + int N = numObjects; + dim3 fullBlocksPerGrid((N + blockSize - 1) / blockSize); + + kernUpdateVelocityBruteForce<<>>(N, dev_pos, dev_vel1, dev_vel2); + kernUpdatePos<<>>(N, dt, dev_pos, dev_vel2); + + cudaMemcpy(dev_vel1, dev_vel2, N * sizeof(glm::vec3), cudaMemcpyDeviceToDevice); } void Boids::stepSimulationScatteredGrid(float dt) { - // TODO-2.1 - // 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. - // - Unstable key sort using Thrust. A stable sort isn't necessary, but you - // are welcome to do a performance comparison. - // - Naively unroll the loop for finding the start and end indices of each - // cell's data pointers in the array of boid indices - // - Perform velocity updates using neighbor search - // - Update positions - // - Ping-pong buffers as needed + + int N = numObjects; + dim3 fullBlocksPerGrid((N + blockSize - 1) / blockSize); + dim3 fullBlocksPerCellGrid((gridCellCount + blockSize - 1) / blockSize); + + kernComputeIndices <<>> (N, gridInverseCellWidth, + dev_pos, dev_particleArrayIndices, dev_particleGridIndices); + checkCUDAErrorWithLine("Error on kernComputeIndices\n"); + + cudaDeviceSynchronize(); + checkCUDAErrorWithLine("Error on sync after computing indices!"); + + sortGridIndices(N, false); + + //TODO: look into making this just a serial operation, rather than parallelizing it inefficiently? + kernFindGridStartEnds <<>> + (N, dev_particleGridIndices, dev_gridCellStartIndices, dev_gridCellEndIndices); + checkCUDAErrorWithLine("Error on kernFindGridStartEnds\n"); + + cudaDeviceSynchronize(); + checkCUDAErrorWithLine("Error on sync after updating gridstarts\n"); + + kernUpdateVelNeighborSearchScattered << > > (N, dev_gridCellStartIndices, dev_gridCellEndIndices, + dev_particleArrayIndices, dev_pos, dev_vel1, dev_vel2); + checkCUDAErrorWithLine("Error on kernUpdateVelNeighborSearchScattered!\n"); + cudaDeviceSynchronize(); + checkCUDAErrorWithLine("Error on sync after updating vel's\n"); + + kernUpdatePos <<>> (N, dt, dev_pos, dev_vel2); + checkCUDAErrorWithLine("Error on kernUpdatePos!\n"); + + cudaMemcpy(dev_vel1, dev_vel2, N * sizeof(glm::vec3), cudaMemcpyDeviceToDevice); + checkCUDAErrorWithLine("cudaMemCpy for ping-pong vel1 and vel2 failed!\n"); + cudaDeviceSynchronize(); } void Boids::stepSimulationCoherentGrid(float dt) { @@ -382,6 +937,45 @@ void Boids::stepSimulationCoherentGrid(float dt) { // - Perform velocity updates using neighbor search // - Update positions // - Ping-pong buffers as needed. THIS MAY BE DIFFERENT FROM BEFORE. + int N = numObjects; + dim3 fullBlocksPerGrid((N + blockSize - 1) / blockSize); + dim3 fullBlocksPerCellGrid((gridCellCount + blockSize - 1) / blockSize); + + kernComputeIndices << > > (N, gridInverseCellWidth, + dev_pos, dev_particleArrayIndices, dev_particleGridIndices); + checkCUDAErrorWithLine("Error on kernComputeIndices\n"); + + cudaDeviceSynchronize(); + checkCUDAErrorWithLine("Error on sync after computing indices!"); + + pvvConverge<<>>(N, dev_pvv, dev_pos, dev_vel1, dev_vel2); + checkCUDAErrorWithLine("Error on sync after converging pvv\n"); + + sortGridIndices(N, true); + + kernFindGridStartEnds <<>> + (N, dev_particleGridIndices, dev_gridCellStartIndices, dev_gridCellEndIndices); + checkCUDAErrorWithLine("Error on kernFindGridStartEnds\n"); + + cudaDeviceSynchronize(); + checkCUDAErrorWithLine("Error on sync after updating gridstarts\n"); + + kernUpdateVelNeighborSearchCoherent <<>> (N, dev_gridCellStartIndices, dev_gridCellEndIndices, + dev_pvv); + checkCUDAErrorWithLine("Error on kernUpdateVelNeighborSearchCoherent\n"); + cudaDeviceSynchronize(); + checkCUDAErrorWithLine("Error on sync after updating vel's\n"); + + pvvSplit <<>> (N, dev_pvv, dev_pos, dev_vel1, dev_vel2); + checkCUDAErrorWithLine("Error on sync after splitting pvv\n"); + + kernUpdatePos << > > (N, dt, dev_pos, dev_vel2); + checkCUDAErrorWithLine("Error on kernUpdatePos!\n"); + + cudaMemcpy(dev_vel1, dev_vel2, N * sizeof(glm::vec3), cudaMemcpyDeviceToDevice); + checkCUDAErrorWithLine("cudaMemCpy for ping-pong vel1 and vel2 failed!\n"); + cudaDeviceSynchronize(); + } void Boids::endSimulation() { @@ -389,6 +983,11 @@ void Boids::endSimulation() { cudaFree(dev_vel2); cudaFree(dev_pos); + cudaFree(dev_particleArrayIndices); + cudaFree(dev_particleGridIndices); + cudaFree(dev_gridCellStartIndices); + cudaFree(dev_gridCellEndIndices); + // TODO-2.1 TODO-2.3 - Free any additional buffers here. } @@ -455,3 +1054,4 @@ void Boids::unitTest() { checkCUDAErrorWithLine("cudaFree failed!"); return; } + diff --git a/src/kernel.h b/src/kernel.h index 3d3da72..687a06f 100644 --- a/src/kernel.h +++ b/src/kernel.h @@ -16,6 +16,8 @@ namespace Boids { void stepSimulationCoherentGrid(float dt); void copyBoidsToVBO(float *vbodptr_positions, float *vbodptr_velocities); + void sortGridIndices(int N, bool coherentGrid); + void endSimulation(); void unitTest(); } diff --git a/src/main.cpp b/src/main.cpp index b82c8c6..2dbb5c2 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; @@ -43,6 +43,15 @@ int main(int argc, char* argv[]) { std::string deviceName; GLFWwindow *window; +//############################## +// CUDA EVENTS FOR TIMING +//############################## +#define TIMEKEEPING_FRAMESIZE 300 +const char timingFileName[] = "../outputData/scratch.csv"; +std::chrono::steady_clock::time_point firstStart, start, stop; +long numSteps = 0; +std::vector eventRecords = {}; + /** * Initialization of CUDA and GLFW. */ @@ -180,6 +189,7 @@ void initShaders(GLuint * program) { } } + //==================================== // Main loop //==================================== @@ -195,6 +205,13 @@ void initShaders(GLuint * program) { cudaGLMapBufferObject((void**)&dptrVertPositions, boidVBO_positions); cudaGLMapBufferObject((void**)&dptrVertVelocities, boidVBO_velocities); + if (numSteps == 0) { + firstStart = std::chrono::high_resolution_clock::now(); + } + if (numSteps % TIMEKEEPING_FRAMESIZE == 0) { + start = std::chrono::high_resolution_clock::now(); + } + // execute the kernel #if UNIFORM_GRID && COHERENT_GRID Boids::stepSimulationCoherentGrid(DT); @@ -204,9 +221,17 @@ void initShaders(GLuint * program) { Boids::stepSimulationNaive(DT); #endif + numSteps++; + #if VISUALIZE Boids::copyBoidsToVBO(dptrVertPositions, dptrVertVelocities); #endif + + if (numSteps % TIMEKEEPING_FRAMESIZE == 0) { + stop = std::chrono::high_resolution_clock::now(); + recordTime(start, stop); + }//if + // unmap buffer object cudaGLUnmapBufferObject(boidVBO_positions); cudaGLUnmapBufferObject(boidVBO_velocities); @@ -256,6 +281,7 @@ void initShaders(GLuint * program) { glfwSwapBuffers(window); #endif } + writeTime(timingFileName); glfwDestroyWindow(window); glfwTerminate(); } @@ -311,3 +337,26 @@ void initShaders(GLuint * program) { glUniformMatrix4fv(location, 1, GL_FALSE, &projection[0][0]); } } + + void recordTime(std::chrono::steady_clock::time_point begin, std::chrono::steady_clock::time_point end) { + long long microseconds = std::chrono::duration_cast(end - begin).count(); + long totalTime = std::chrono::duration_cast(end - firstStart).count(); + double numMs = microseconds / 1000; + + eventRecords.push_back({ numSteps, numMs, totalTime }); + + }//recordTime + + void writeTime(const char* fileName) { + FILE* of = fopen(fileName, "w"); + + for (auto& record : eventRecords) { + double millisPerFrame = record.time / TIMEKEEPING_FRAMESIZE; + millisPerFrame /= 1000.0;//seconds per frame + double fps = 1.0 / millisPerFrame; + double seconds = record.totalTime / 1000.0; + fprintf(of, "%d,%0.3f,%f\n", record.frameNo, seconds, fps); + }//for + + fclose(of); + }//writeTime diff --git a/src/main.hpp b/src/main.hpp index 88e9df7..e972b11 100644 --- a/src/main.hpp +++ b/src/main.hpp @@ -14,6 +14,15 @@ #include "utilityCore.hpp" #include "glslUtility.hpp" #include "kernel.h" +#include +#include +#include + +typedef struct timeRecord { + int frameNo; + double time;//milliseconds + long totalTime;//milliseconds +} timeRecord; //==================================== // GL Stuff @@ -78,3 +87,6 @@ void runCUDA(); bool init(int argc, char **argv); void initVAO(); void initShaders(GLuint *program); + +void recordTime(std::chrono::steady_clock::time_point begin, std::chrono::steady_clock::time_point end); +void writeTime(const char* fileName);