Project 2: Yilin Liu

README.md

@@ -3,12 +3,108 @@ 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)
+* Yilin Liu
+  * [LinkedIn](https://www.linkedin.com/in/yilin-liu-9538ba1a5/)
+  * [Personal website](https://www.yilin.games)
+* Tested on personal laptop:
+  - Windows 10, Intel(R) Core(TM), i7-10750H CPU @ 2.60GHz 2.59 GHz, RTX 2070 Max-Q 8GB
 
-### (TODO: Your README)
+## Features
+This project uses CUDA to implement and improve a number of parallelized scan and stream compaction algorithms. Following features have been implemented:
 
-Include analysis, etc. (Remember, this is public, so don't put
-anything here that you don't want to share with the world.)
+- Scan: calculate the prefix sum (using arbitrary operator) of an array
 
+  - CPU scan with/without simulating parallelized scan 
+
+  - GPU naive scan 
+
+  - GPU work-efficient scan 
+
+  - GPU scan using `thrust::exclusive_scan` 
+
+- Stream compaction: remove elements that unmeet specific conditions from an array, and keep the rest compact in memory
+
+  - CPU stream compaction with/without CPU scan
+
+  - GPU stream compaction using the work-efficient scan
+
+
+
+
+## Reflection
+
+* Based on array size of 33554432 (2^25), I optimized each method by selecting the block size which could provide best performance.
+
+  | Methods  | Optimized Block Size |
+  | ------------- | ------------- |
+  | Naive Scan  | 128  |
+  | Work Efficient  | 512  |
+  | Thrust  | 512  |
+
+
+* Based on the optimized block size, I compare all of these GPU Scan implementations (Naive, Work-Efficient, and Thrust) to the serial CPU version of scan. 
+
+  |![image](./img/table1.png)|
+  |:--:| 
+  | *Comparision of Scan Methods vs Time* |
+* From the figure above, we can see that there is no significant differences between CPU, Naive and Work Efficient method while the array size is less than 15M, although Work Efficient method is slightly faster than the other two. However, as the array size increases, work efficient method is much faster while other two continue to grow linearly. On the other side, the Thrust method has always been the fastest method and take almost constant time as array increases. 
+* The naive and work efficient method are slow since GPU has to read from global memory, which is very costly. The reason why Thrust library is super fast could probably be that it saves copy/paste cost among kernels.
+
+  |![image](./img/table2.png)|
+  |:--:| 
+  | *Comparision of Compaction Methods vs Time* |
+* For the stream compaction algorithm, the GPU method surpassed all CPU mthods. The stream compaction method using scan is the slowest one since it has to do more operations from global memory. 
+
+
+
+
+## Example Output for Array Size of 2^25
+```
+
+****************
+** SCAN TESTS **
+****************
+    [  24  27   7  15  47  21  25  25   4  30  41  18  28 ...  11   0 ]
+==== cpu scan, power-of-two ====
+   elapsed time: 49.3806ms    (std::chrono Measured)
+==== cpu scan, non-power-of-two ====
+   elapsed time: 48.9979ms    (std::chrono Measured)
+    passed
+==== naive scan, power-of-two ====
+   elapsed time: 52.6029ms    (CUDA Measured)
+    passed
+==== naive scan, non-power-of-two ====
+   elapsed time: 50.8384ms    (CUDA Measured)
+    passed
+==== work-efficient scan, power-of-two ====
+   elapsed time: 21.571ms    (CUDA Measured)
+    passed
+==== work-efficient scan, non-power-of-two ====
+   elapsed time: 22.2183ms    (CUDA Measured)
+    passed
+==== thrust scan, power-of-two ====
+   elapsed time: 1.02707ms    (CUDA Measured)
+    passed
+==== thrust scan, non-power-of-two ====
+   elapsed time: 1.06106ms    (CUDA Measured)
+    passed
+
+*****************************
+** STREAM COMPACTION TESTS **
+*****************************
+    [   2   2   3   3   2   2   0   1   3   0   2   2   0 ...   2   0 ]
+==== cpu compact without scan, power-of-two ====
+   elapsed time: 74.2421ms    (std::chrono Measured)
+    passed
+==== cpu compact without scan, non-power-of-two ====
+   elapsed time: 74.3851ms    (std::chrono Measured)
+    passed
+==== cpu compact with scan ====
+   elapsed time: 166.162ms    (std::chrono Measured)
+    passed
+==== work-efficient compact, power-of-two ====
+   elapsed time: 26.1186ms    (CUDA Measured)
+    passed
+==== work-efficient compact, non-power-of-two ====
+   elapsed time: 26.5175ms    (CUDA Measured)
+```

src/main.cpp

@@ -5,6 +5,7 @@
  * @date      2015
  * @copyright University of Pennsylvania
  */
+#include <vector>
 
 #include <cstdio>
 #include <stream_compaction/cpu.h>
@@ -13,15 +14,14 @@
 #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 << 26; // 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 main(int argc, char* argv[]) {
+void scanTest() {
     // Scan tests
-
+    std::vector<float> timeVector;
     printf("\n");
     printf("****************\n");
     printf("** SCAN TESTS **\n");
@@ -37,20 +37,20 @@ int main(int argc, char* argv[]) {
     zeroArray(SIZE, b);
     printDesc("cpu scan, power-of-two");
     StreamCompaction::CPU::scan(SIZE, b, a);
-    printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)");
-    printArray(SIZE, b, true);
+    printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), timeVector, "(std::chrono Measured)");
+    //printArray(SIZE, b, true);
 
     zeroArray(SIZE, c);
     printDesc("cpu scan, non-power-of-two");
     StreamCompaction::CPU::scan(NPOT, c, a);
-    printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)");
-    printArray(NPOT, b, true);
+    printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), timeVector, "(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)");
+    printElapsedTime(StreamCompaction::Naive::timer().getGpuElapsedTimeForPreviousOperation(), timeVector, "(CUDA Measured)");
     //printArray(SIZE, c, true);
     printCmpResult(SIZE, b, c);
 
@@ -63,35 +63,35 @@ int main(int argc, char* argv[]) {
     zeroArray(SIZE, c);
     printDesc("naive scan, non-power-of-two");
     StreamCompaction::Naive::scan(NPOT, c, a);
-    printElapsedTime(StreamCompaction::Naive::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)");
+    printElapsedTime(StreamCompaction::Naive::timer().getGpuElapsedTimeForPreviousOperation(), timeVector, "(CUDA Measured)");
     //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)");
+    printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), timeVector, "(CUDA Measured)");
     //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)");
+    printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), timeVector, "(CUDA Measured)");
     //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)");
+    printElapsedTime(StreamCompaction::Thrust::timer().getGpuElapsedTimeForPreviousOperation(), timeVector, "(CUDA Measured)");
     //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)");
+    printElapsedTime(StreamCompaction::Thrust::timer().getGpuElapsedTimeForPreviousOperation(), timeVector, "(CUDA Measured)");
     //printArray(NPOT, c, true);
     printCmpResult(NPOT, b, c);
 
@@ -113,42 +113,51 @@ int main(int argc, char* argv[]) {
     zeroArray(SIZE, b);
     printDesc("cpu compact without scan, power-of-two");
     count = StreamCompaction::CPU::compactWithoutScan(SIZE, b, a);
-    printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)");
+    printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), timeVector, "(std::chrono Measured)");
     expectedCount = count;
-    printArray(count, b, true);
+    //printArray(count, b, true);
     printCmpLenResult(count, expectedCount, b, b);
 
     zeroArray(SIZE, c);
     printDesc("cpu compact without scan, non-power-of-two");
     count = StreamCompaction::CPU::compactWithoutScan(NPOT, c, a);
-    printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)");
+    printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), timeVector, "(std::chrono Measured)");
     expectedNPOT = count;
-    printArray(count, c, true);
+    //printArray(count, c, true);
     printCmpLenResult(count, expectedNPOT, b, c);
 
     zeroArray(SIZE, c);
     printDesc("cpu compact with scan");
     count = StreamCompaction::CPU::compactWithScan(SIZE, c, a);
-    printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)");
-    printArray(count, c, true);
+    printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), timeVector, "(std::chrono Measured)");
+    //printArray(count, c, true);
     printCmpLenResult(count, expectedCount, b, c);
 
     zeroArray(SIZE, c);
     printDesc("work-efficient compact, power-of-two");
     count = StreamCompaction::Efficient::compact(SIZE, c, a);
-    printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)");
+    printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), timeVector, "(CUDA Measured)");
     //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)");
+    printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), timeVector, "(CUDA Measured)");
     //printArray(count, c, true);
     printCmpLenResult(count, expectedNPOT, b, c);
 
+
+    for (int i = 0; i < timeVector.size(); i++)
+    {
+        std::cout << timeVector[i] << std::endl;
+    }
     system("pause"); // stop Win32 console from closing on exit
     delete[] a;
     delete[] b;
     delete[] c;
 }
+
+int main(int argc, char* argv[]) {
+    scanTest();
+}

src/testing_helpers.hpp

@@ -5,7 +5,7 @@
 #include <iostream>
 #include <string>
 #include <ctime>
-
+#include <vector>
 template<typename T>
 int cmpArrays(int n, T *a, T *b) {
     for (int i = 0; i < n; i++) {
@@ -68,9 +68,10 @@ void printArray(int n, int *a, bool abridged = false) {
     }
     printf("]\n");
 }
-
-template<typename T>
-void printElapsedTime(T time, std::string note = "")
+//
+//template<typename T, typename A>
+void printElapsedTime(float time, std::vector<float>& timeArray, std::string note = "")
 {
     std::cout << "   elapsed time: " << time << "ms    " << note << std::endl;
+    timeArray.push_back(time);
 }

stream_compaction/common.cu

@@ -24,6 +24,11 @@ namespace StreamCompaction {
          */
         __global__ void kernMapToBoolean(int n, int *bools, const int *idata) {
             // TODO
+            int index = threadIdx.x + (blockIdx.x * blockDim.x);
+            if (index >= n) {
+                return;
+            }
+            bools[index] = idata[index] ? 1 : 0;
         }
 
         /**
@@ -33,6 +38,13 @@ namespace StreamCompaction {
         __global__ void kernScatter(int n, int *odata,
                 const int *idata, const int *bools, const int *indices) {
             // TODO
+            int index = threadIdx.x + (blockIdx.x * blockDim.x);
+            if (index >= n) {
+                return;
+            }
+            if (bools[index]) {
+                odata[indices[index]] = idata[index];
+            }
         }
 
     }

stream_compaction/common.h

@@ -13,6 +13,8 @@
 #define FILENAME (strrchr(__FILE__, '/') ? strrchr(__FILE__, '/') + 1 : __FILE__)
 #define checkCUDAError(msg) checkCUDAErrorFn(msg, FILENAME, __LINE__)
 
+#define BLOCK_SIZE 128
+
 /**
  * Check for CUDA errors; print and exit if there was a problem.
  */

stream_compaction/cpu.cu

@@ -20,9 +20,20 @@ namespace StreamCompaction {
         void scan(int n, int *odata, const int *idata) {
             timer().startCpuTimer();
             // TODO
+            odata[0] = 0;
+            odata[1] = idata[0];
+            for (size_t i = 2; i < n; i++) {
+                odata[i] = odata[i - 1] + idata[i - 1];
+            }
             timer().endCpuTimer();
         }
-
+        void scanInclusive(int n, int* odata, const int* idata) {
+            // TODO
+            odata[0] = idata[0];
+            for (size_t i = 1; i < n; i++) {
+                odata[i] = odata[i - 1] + idata[i];
+            }
+        }
         /**
          * CPU stream compaction without using the scan function.
          *
@@ -31,8 +42,15 @@ namespace StreamCompaction {
         int compactWithoutScan(int n, int *odata, const int *idata) {
             timer().startCpuTimer();
             // TODO
+            int j = 0;
+            for (size_t i = 0; i < n; i++) {
+                if (idata[i] != 0) {
+                    odata[j] = idata[i];
+                    j++;
+                }
+            }
             timer().endCpuTimer();
-            return -1;
+            return j;
         }
 
         /**
@@ -41,10 +59,25 @@ namespace StreamCompaction {
          * @returns the number of elements remaining after compaction.
          */
         int compactWithScan(int n, int *odata, const int *idata) {
+            int* boolFlag = new int[n]; // temporary array 
+            int* scanRes = new int[n];
             timer().startCpuTimer();
             // TODO
+            for (size_t i = 0; i < n; i++)
+            {
+                boolFlag[i] = idata[i] == 0 ? 0 : 1;
+            }
+            scanInclusive(n, scanRes, boolFlag); // odata: scan result
+
+
+            for (size_t i = 0; i < n; i++) {
+                if (boolFlag[i] == 0) continue;
+                odata[scanRes[i] - 1] = idata[i];
+            }
+
             timer().endCpuTimer();
-            return -1;
+            return scanRes[n - 1];
+
         }
     }
 }