Project 2: Jacky Lu

README.md

@@ -3,12 +3,80 @@ 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)
+* Jacky Lu
+  * [LinkedIn](https://www.linkedin.com/in/jacky-lu-506968129/)
 
-### (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.)
+# README
+
+## Result:
+### The implementations' block sizes seem to be optimized at 128 threads per block.
+
+### Performance Comparison Between GPU Scan Implementations (Naive, Work-Efficient, and Thrust) And CPU Scan Implementations
+#### (Tested With 128 Threads Per Block)
+![](img/plot.png)
+
+### Output Of The Test Program With 33,554,432 Array Elements (Tested With 128 Threads Per Block)
+```
+****************
+** SCAN TESTS **
+****************
+    [  12   3  19  17  37  34   6  17  18  22  23   8  49 ...  17   0 ]
+==== cpu scan, power-of-two ====
+   elapsed time: 102.834ms    (std::chrono Measured)
+    [   0  12  15  34  51  88 122 128 145 163 185 208 216 ... 821676053 821676070 ]
+==== cpu scan, non-power-of-two ====
+   elapsed time: 38.8949ms    (std::chrono Measured)
+    [   0  12  15  34  51  88 122 128 145 163 185 208 216 ... 821675987 821676034 ]
+    passed
+==== naive scan, power-of-two ====
+   elapsed time: 19.5621ms    (CUDA Measured)
+    [   0  12  15  34  51  88 122 128 145 163 185 208 216 ... 821676053 821676070 ]
+    passed
+==== naive scan, non-power-of-two ====
+   elapsed time: 19.9352ms    (CUDA Measured)
+    [   0  12  15  34  51  88 122 128 145 163 185 208 216 ... 821675987 821676034 ]
+    passed
+==== work-efficient scan, power-of-two ====
+   elapsed time: 14.7942ms    (CUDA Measured)
+    [   0  12  15  34  51  88 122 128 145 163 185 208 216 ... 821676053 821676070 ]
+    passed
+==== work-efficient scan, non-power-of-two ====
+   elapsed time: 14.8075ms    (CUDA Measured)
+    [   0  12  15  34  51  88 122 128 145 163 185 208 216 ... 821675987 821676034 ]
+    passed
+==== thrust scan, power-of-two ====
+   elapsed time: 0.8088ms    (CUDA Measured)
+    [   0  12  15  34  51  88 122 128 145 163 185 208 216 ... 821676053 821676070 ]
+    passed
+==== thrust scan, non-power-of-two ====
+   elapsed time: 0.806624ms    (CUDA Measured)
+    [   0  12  15  34  51  88 122 128 145 163 185 208 216 ... 821675987 821676034 ]
+    passed
+
+*****************************
+** STREAM COMPACTION TESTS **
+*****************************
+    [   0   2   1   2   0   0   3   3   1   2   2   2   2 ...   2   0 ]
+==== cpu compact without scan, power-of-two ====
+   elapsed time: 60.1642ms    (std::chrono Measured)
+    [   2   1   2   3   3   1   2   2   2   2   3   1   3 ...   3   2 ]
+    passed
+==== cpu compact without scan, non-power-of-two ====
+   elapsed time: 60.4809ms    (std::chrono Measured)
+    [   2   1   2   3   3   1   2   2   2   2   3   1   3 ...   1   3 ]
+    passed
+==== cpu compact with scan ====
+   elapsed time: 205.018ms    (std::chrono Measured)
+    [   2   1   2   3   3   1   2   2   2   2   3   1   3 ...   3   2 ]
+    passed
+==== work-efficient compact, power-of-two ====
+   elapsed time: 17.7603ms    (CUDA Measured)
+    [   2   1   2   3   3   1   2   2   2   2   3   1   3 ...   3   2 ]
+    passed
+==== work-efficient compact, non-power-of-two ====
+   elapsed time: 17.4516ms    (CUDA Measured)
+    [   2   1   2   3   3   1   2   2   2   2   3   1   3 ...   1   3 ]
+    passed
+```
 

src/main.cpp

@@ -13,7 +13,7 @@
 #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 << 25); // feel free to change the size of array (8 originally)
 const int NPOT = SIZE - 3; // Non-Power-Of-Two
 int *a = new int[SIZE];
 int *b = new int[SIZE];
@@ -51,7 +51,7 @@ int main(int argc, char* argv[]) {
     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
@@ -64,35 +64,35 @@ int main(int argc, char* argv[]) {
     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(NPOT, 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 +137,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

stream_compaction/common.cu

@@ -1,5 +1,7 @@
 #include "common.h"
 
+#include <device_launch_parameters.h>
+
 void checkCUDAErrorFn(const char *msg, const char *file, int line) {
     cudaError_t err = cudaGetLastError();
     if (cudaSuccess == err) {
@@ -24,6 +26,15 @@ namespace StreamCompaction {
          */
         __global__ void kernMapToBoolean(int n, int *bools, const int *idata) {
             // TODO
+            int index = (blockIdx.x * blockDim.x) + threadIdx.x;
+            if (index < n) {
+                if (idata[index] != 0) {
+                    bools[index] = 1;
+                } else {
+                    bools[index] = 0;
+                }
+
+            }
         }
 
         /**
@@ -33,6 +44,10 @@ namespace StreamCompaction {
         __global__ void kernScatter(int n, int *odata,
                 const int *idata, const int *bools, const int *indices) {
             // TODO
+            int index = (blockIdx.x * blockDim.x) + threadIdx.x;
+            if (index < n && bools[index] == 1) {
+                odata[indices[index]] = idata[index];
+            }
         }
 
     }

stream_compaction/cpu.cu

@@ -20,6 +20,10 @@ namespace StreamCompaction {
         void scan(int n, int *odata, const int *idata) {
             timer().startCpuTimer();
             // TODO
+            odata[0] = 0;
+            for (int i = 1; i < n; i++) {
+                odata[i] = odata[i - 1] + idata[i - 1];
+            }
             timer().endCpuTimer();
         }
 
@@ -31,8 +35,15 @@ namespace StreamCompaction {
         int compactWithoutScan(int n, int *odata, const int *idata) {
             timer().startCpuTimer();
             // TODO
+            int counter = 0;
+            for (int i = 0; i < n; i++) {
+                if (idata[i] != 0) {
+                    odata[counter] = idata[i];
+                    counter++;
+                }
+            }
             timer().endCpuTimer();
-            return -1;
+            return counter;
         }
 
         /**
@@ -43,8 +54,32 @@ namespace StreamCompaction {
         int compactWithScan(int n, int *odata, const int *idata) {
             timer().startCpuTimer();
             // TODO
+            // Construct criteria vector
+            int* criteriaVec = new int[n];
+            for (int i = 0; i < n; i++) {
+                criteriaVec[i] = idata[i] != 0 ? 1 : 0;
+            }
+
+            // Construct scan vector from criteria vector
+            int* scanVec = new int[n];
+            scanVec[0] = 0;
+            for (int i = 1; i < n; i++) {
+                scanVec[i] = scanVec[i - 1] + criteriaVec[i - 1];
+            }
+
+            // Scatter
+            int counter = 0;
+            for (int i = 0; i < n; i++) {
+                if (criteriaVec[i] == 1) {
+                    odata[scanVec[i]] = idata[i];
+                    counter++;
+                }
+            }
+
+            delete[] criteriaVec;
+            delete[] scanVec;
             timer().endCpuTimer();
-            return -1;
+            return counter;
         }
     }
 }