Project 2: Liam Dugan

README.md

@@ -3,12 +3,106 @@ CUDA Stream Compaction
 
 **University of Pennsylvania, CIS 565: GPU Programming and Architecture, Project 2**
 
-* (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)
+* Liam Dugan
+  * [LinkedIn](https://www.linkedin.com/in/liam-dugan-95a961135/), [personal website](http://liamdugan.com/)
+* Tested on: Windows 10, Ryzen 5 1600 @ 3.20GHz 16GB, GTX 1070 16GB (Personal Computer)
 
-### (TODO: Your README)
 
-Include analysis, etc. (Remember, this is public, so don't put
-anything here that you don't want to share with the world.)
+# Stream Compaction
+This homework is an introduction to implementing and optimizing parallel algorithms on the GPU. To do this we were tasked with writing a CPU implementation and various GPU implementations and comparing their timing results.
 
+## Features Completed
+* CPU Scan and stream compaction implementation
+* GPU Naive scan implementation
+* GPU Work efficient scan and stream compaction implementation
+* Wrapper for Thrust compaction
+
+# Questions
+### Roughly optimize the block sizes of each of your implementations for minimal run time on your GPU
+I tested block sizes from 32 up to 1024 in power of two increments and determined that a block size of 64 was ideal for the naive implementation and a block size of 128 was ideal for the work-efficient implementations.
+
+For the timing results, an array of size 2^20 (~1 million) was tested while changing the block size. Time is in miliseconds.
+
+![](images/BlockSize.png)
+
+### Compare all of these GPU Scan implementations (Naive, Work-Efficient, and Thrust) to the serial CPU version of Scan. Plot a graph of the comparison (with array size on the independent axis).
+
+![](images/Scan.png)
+
+![](images/StreamCompaction.png)
+
+### Write a brief explanation of the phenomena you see here.
+I believe the main performance bottleneck in my scan code is undoubtedly the memory accesses. In both my naive and work efficient scan implementations I write to global memory after every level of the tree is traversed (i.e. log(n) times). This is significantly slower than using shared memory, which would only write to global memory log(n / blocksize) times.
+
+Additionally, in my work efficient scan algorithm I use ping-pong buffers,but memcpy the data between buffers on every cycle of both the upsweep and the downsweep of the algorithm. Since I do not have to do this in my naive implementation, I believe that is why my work efficient scan is slower.
+
+![](images/Code.PNG)
+
+For the thrust scan, the reason why I believe the power of two array test runs so much slower than the non power of two code is due to some sort of internal thrust library bookkeeping. I believe once a thrust function is called, there is some sort of one-time-only process to initialize thrust specific state. Thus when we call the power-of-two length array thrust scan it has to take that extra set up time, but once we call the non power of two code, all the setup has already been completed and it can run quickly.
+
+
+### Test Program Output
+
+```
+
+****************
+** SCAN TESTS **
+****************
+    [  12   6  12  29  23  44  41  42  46  40   9  23  22 ...   6   0 ]
+==== cpu scan, power-of-two ====
+   elapsed time: 1.47776ms    (std::chrono Measured)
+    [   0  12  18  30  59  82 126 167 209 255 295 304 327 ... 51356118 51356124 ]
+==== cpu scan, non-power-of-two ====
+   elapsed time: 1.57856ms    (std::chrono Measured)
+    [   0  12  18  30  59  82 126 167 209 255 295 304 327 ... 51356052 51356083 ]
+    passed
+==== naive scan, power-of-two ====
+   elapsed time: 5.05549ms    (CUDA Measured)
+    [   0  12  18  30  59  82 126 167 209 255 295 304 327 ... 51356118 51356124 ]
+    passed
+==== naive scan, non-power-of-two ====
+   elapsed time: 5.13843ms    (CUDA Measured)
+    [   0  12  18  30  59  82 126 167 209 255 295 304 327 ... 51356118 51356124 ]
+    passed
+==== work-efficient scan, power-of-two ====
+   elapsed time: 9.00813ms    (CUDA Measured)
+    [   0  12  18  30  59  82 126 167 209 255 295 304 327 ... 51356118 51356124 ]
+    passed
+==== work-efficient scan, non-power-of-two ====
+   elapsed time: 8.82483ms    (CUDA Measured)
+    [   0  12  18  30  59  82 126 167 209 255 295 304 327 ... 51356052 51356083 ]
+    passed
+==== thrust scan, power-of-two ====
+   elapsed time: 4.64384ms    (CUDA Measured)
+    [   0  12  18  30  59  82 126 167 209 255 295 304 327 ... 51356118 51356124 ]
+    passed
+==== thrust scan, non-power-of-two ====
+   elapsed time: 0.334848ms    (CUDA Measured)
+    [   0  12  18  30  59  82 126 167 209 255 295 304 327 ... 51356052 51356083 ]
+    passed
+
+*****************************
+** STREAM COMPACTION TESTS **
+*****************************
+    [   3   1   3   3   0   1   1   1   0   2   1   1   1 ...   3   0 ]
+==== cpu compact without scan, power-of-two ====
+   elapsed time: 5.27103ms    (std::chrono Measured)
+    [   3   1   3   3   1   1   1   2   1   1   1   2   3 ...   3   3 ]
+    passed
+==== cpu compact without scan, non-power-of-two ====
+   elapsed time: 5.21855ms    (std::chrono Measured)
+    [   3   1   3   3   1   1   1   2   1   1   1   2   3 ...   3   3 ]
+    passed
+==== cpu compact with scan ====
+   elapsed time: 12.4281ms    (std::chrono Measured)
+    [   3   1   3   3   1   1   1   2   1   1   1   2   3 ...   3   3 ]
+    passed
+==== work-efficient compact, power-of-two ====
+   elapsed time: 7.15469ms    (CUDA Measured)
+    [   3   1   3   3   1   1   1   2   1   1   1   2   3 ...   3   3 ]
+    passed
+==== work-efficient compact, non-power-of-two ====
+   elapsed time: 7.18848ms    (CUDA Measured)
+    [   3   1   3   3   1   1   1   2   1   1   1   2   3 ...   3   3 ]
+    passed
+```

src/main.cpp

@@ -13,11 +13,12 @@
 #include <stream_compaction/thrust.h>
 #include "testing_helpers.hpp"
 
-const int SIZE = 1 << 8; // feel free to change the size of array
+const int SIZE = 1 << 21; // feel free to change the size of array
 const int NPOT = SIZE - 3; // Non-Power-Of-Two
 int *a = new int[SIZE];
 int *b = new int[SIZE];
 int *c = new int[SIZE];
+int *d = new int[SIZE];
 
 int main(int argc, char* argv[]) {
     // Scan tests
@@ -46,53 +47,53 @@ int main(int argc, char* argv[]) {
     printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)");
     printArray(NPOT, b, true);
     printCmpResult(NPOT, b, c);
-
+    
     zeroArray(SIZE, c);
     printDesc("naive scan, power-of-two");
     StreamCompaction::Naive::scan(SIZE, c, a);
     printElapsedTime(StreamCompaction::Naive::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)");
-    //printArray(SIZE, c, true);
+    printArray(SIZE, c, true);
     printCmpResult(SIZE, b, c);
 
-	/* For bug-finding only: Array of 1s to help find bugs in stream compaction or scan
-	onesArray(SIZE, c);
-	printDesc("1s array for finding bugs");
-	StreamCompaction::Naive::scan(SIZE, c, a);
-	printArray(SIZE, c, true); */
-
+    /* For bug-finding only: Array of 1s to help find bugs in stream compaction or scan */
+/*    onesArray(SIZE, d);
+    printDesc("1s array for finding bugs");
+    StreamCompaction::Efficient::scan(SIZE, c, d);
+    printArray(SIZE, c, true);
+    */
     zeroArray(SIZE, c);
     printDesc("naive scan, non-power-of-two");
     StreamCompaction::Naive::scan(NPOT, c, a);
     printElapsedTime(StreamCompaction::Naive::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)");
-    //printArray(SIZE, c, true);
+    printArray(SIZE, c, true);
     printCmpResult(NPOT, b, c);
 
     zeroArray(SIZE, c);
     printDesc("work-efficient scan, power-of-two");
     StreamCompaction::Efficient::scan(SIZE, c, a);
     printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)");
-    //printArray(SIZE, c, true);
+    printArray(SIZE, c, true);
     printCmpResult(SIZE, b, c);
 
     zeroArray(SIZE, c);
     printDesc("work-efficient scan, non-power-of-two");
     StreamCompaction::Efficient::scan(NPOT, c, a);
     printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)");
-    //printArray(NPOT, c, true);
+    printArray(NPOT, c, true);
     printCmpResult(NPOT, b, c);
 
     zeroArray(SIZE, c);
     printDesc("thrust scan, power-of-two");
     StreamCompaction::Thrust::scan(SIZE, c, a);
     printElapsedTime(StreamCompaction::Thrust::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)");
-    //printArray(SIZE, c, true);
+    printArray(SIZE, c, true);
     printCmpResult(SIZE, b, c);
 
     zeroArray(SIZE, c);
     printDesc("thrust scan, non-power-of-two");
     StreamCompaction::Thrust::scan(NPOT, c, a);
     printElapsedTime(StreamCompaction::Thrust::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)");
-    //printArray(NPOT, c, true);
+    printArray(NPOT, c, true);
     printCmpResult(NPOT, b, c);
 
     printf("\n");
@@ -137,14 +138,14 @@ int main(int argc, char* argv[]) {
     printDesc("work-efficient compact, power-of-two");
     count = StreamCompaction::Efficient::compact(SIZE, c, a);
     printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)");
-    //printArray(count, c, true);
+    printArray(count, c, true);
     printCmpLenResult(count, expectedCount, b, c);
 
     zeroArray(SIZE, c);
     printDesc("work-efficient compact, non-power-of-two");
     count = StreamCompaction::Efficient::compact(NPOT, c, a);
     printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)");
-    //printArray(count, c, true);
+    printArray(count, c, true);
     printCmpLenResult(count, expectedNPOT, b, c);
 
     system("pause"); // stop Win32 console from closing on exit

src/testing_helpers.hpp

@@ -1,8 +1,8 @@
 #pragma once
 
-#include <cstdlib>
-#include <cstdio>
-#include <iostream>
+#include <cstdlib>
+#include <cstdio>
+#include <iostream>
 #include <string>
 #include <ctime>
 
@@ -69,8 +69,8 @@ void printArray(int n, int *a, bool abridged = false) {
     printf("]\n");
 }
 
-template<typename T>
-void printElapsedTime(T time, std::string note = "")
-{
-	std::cout << "   elapsed time: " << time << "ms    " << note << std::endl;
+template<typename T>
+void printElapsedTime(T time, std::string note = "")
+{
+	std::cout << "   elapsed time: " << time << "ms    " << note << std::endl;
 }

stream_compaction/CMakeLists.txt

@@ -13,5 +13,5 @@ set(SOURCE_FILES
 
 cuda_add_library(stream_compaction
     ${SOURCE_FILES}
-    OPTIONS -arch=sm_20
+    OPTIONS -arch=sm_60
     )

stream_compaction/common.cu

@@ -22,8 +22,16 @@ namespace StreamCompaction {
          * Maps an array to an array of 0s and 1s for stream compaction. Elements
          * which map to 0 will be removed, and elements which map to 1 will be kept.
          */
-        __global__ void kernMapToBoolean(int n, int *bools, const int *idata) {
-            // TODO
+        __global__ void kernMapToBoolean(int n, int paddedN, int *bools, const int *idata) {
+          // get index first and reject if greater than paddedN
+          int index = threadIdx.x + (blockIdx.x * blockDim.x);
+          if (index >= paddedN)
+          {
+            return;
+          }
+
+          // determine if you're a boolean (if you're in the part that's just padded on, give yourself a 0)
+          bools[index] = (idata[index] && index < n) ? 1 : 0;
         }
 
         /**
@@ -32,8 +40,18 @@ namespace StreamCompaction {
          */
         __global__ void kernScatter(int n, int *odata,
                 const int *idata, const int *bools, const int *indices) {
-            // TODO
-        }
+
+          // get index first
+          int index = threadIdx.x + (blockIdx.x * blockDim.x);
+          if (index >= n)
+          {
+            return;
+          }
 
+          if (bools[index])
+          {
+            odata[indices[index]] = idata[index];
+          }
+        }
     }
 }

stream_compaction/common.h

@@ -32,7 +32,7 @@ inline int ilog2ceil(int x) {
 
 namespace StreamCompaction {
     namespace Common {
-        __global__ void kernMapToBoolean(int n, int *bools, const int *idata);
+        __global__ void kernMapToBoolean(int n, int paddedN, int *bools, const int *idata);
 
         __global__ void kernScatter(int n, int *odata,
                 const int *idata, const int *bools, const int *indices);

stream_compaction/cpu.cu

@@ -1,15 +1,15 @@
 #include <cstdio>
 #include "cpu.h"
 
-#include "common.h"
+#include "common.h"
 
 namespace StreamCompaction {
     namespace CPU {
-        using StreamCompaction::Common::PerformanceTimer;
-        PerformanceTimer& timer()
-        {
-	        static PerformanceTimer timer;
-	        return timer;
+        using StreamCompaction::Common::PerformanceTimer;
+        PerformanceTimer& timer()
+        {
+	        static PerformanceTimer timer;
+	        return timer;
         }
 
         /**
@@ -19,7 +19,14 @@ namespace StreamCompaction {
          */
         void scan(int n, int *odata, const int *idata) {
 	        timer().startCpuTimer();
-            // TODO
+
+          int sum = 0;
+          for (int i = 0; i < n; ++i)
+          {
+            odata[i] = sum;
+            sum += idata[i];
+          }
+
 	        timer().endCpuTimer();
         }
 
@@ -30,9 +37,20 @@ namespace StreamCompaction {
          */
         int compactWithoutScan(int n, int *odata, const int *idata) {
 	        timer().startCpuTimer();
-            // TODO
+
+          int index = 0;
+          for (int i = 0; i < n; ++i)
+          {
+            // if the data meets the condition put it in
+            if (idata[i])
+            {
+              odata[index] = idata[i];
+              ++index;
+            }
+          }
+
 	        timer().endCpuTimer();
-            return -1;
+          return index;
         }
 
         /**
@@ -41,10 +59,33 @@ namespace StreamCompaction {
          * @returns the number of elements remaining after compaction.
          */
         int compactWithScan(int n, int *odata, const int *idata) {
-	        timer().startCpuTimer();
-	        // TODO
+          int* scanned = (int*) malloc(sizeof(int) * n);
+
+          timer().startCpuTimer();
+
+          int sum = 0;
+          for (int i = 0; i < n; ++i)
+          {
+            scanned[i] = sum;
+            if (idata[i]) 
+            {
+              ++sum;
+            }
+          }
+
+          // now scatter
+          for (int j = 0; j < n; j++)
+          {
+            if (idata[j])
+            {
+              odata[scanned[j]] = idata[j];
+            }
+          }
+
 	        timer().endCpuTimer();
-            return -1;
+
+          free(scanned);
+          return sum;
         }
     }
 }