Project 2: Zirui Zang

README.md

@@ -3,12 +3,41 @@ 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)
+* Zirui Zang
+  * [LinkedIn](https://www.linkedin.com/in/zirui-zang/)
+* Tested on: Windows 10, AMD Ryzen 7 3700X @ 3.60GHz 32GB, RTX2070 SUPER 8GB (Personal)
+* I'm using one late day for this project. 
 
-### (TODO: Your README)
+### Results
 
-Include analysis, etc. (Remember, this is public, so don't put
-anything here that you don't want to share with the world.)
+In this lab I have implemented the scan (prefix-sum) and stream compaction alogrithm in three different ways:
 
+1. CPU sequential method
+2. GPU naive method, which parallelized the addition but has to iterate log(n) times.
+3. GPU efficient method, which parallelized the whole process.
+
+Additional attention has been given to utilizing shared memory in kernel computation to reduce global memory access. The array that needs to be operate on is pre-fetch into shared memory to scan operations can run within shared memory.
+
+The code is also implementation so large input array can be split into GPU blocksizes and merged after computation.
+
+Here are the charts of my implementations:
+![chart_scan](img/chart_scan.png)
+
+![chart_stream](img/chart_stream.png)
+
+Here are the results of my implementations:
+![scan](img/scan.png)
+
+![string](img/string.png)
+
+Here are some issues with the code:
+1. The efficient method is not efficient at all.
+
+### Discussions
+
+The efficient method is not so efficient mainly due to complex kernel operations and excessive memory operations. In the upper and down sweep more than half of the threads are idle. This creates a lot of CUDA core occupation and very scattering memory access. If the numbers in the array can be more comgregated, we can launch fewer threads and save more time in seqential memory access.
+
+
+### VS Studio useless machine
+![VS studio](img/vss.gif)
+This is probably due to the autosync of my cloud.

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 << 18; // 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];
@@ -26,8 +26,10 @@ int main(int argc, char* argv[]) {
     printf("****************\n");
     printf("** SCAN TESTS **\n");
     printf("****************\n");
+    printf("%d ", SIZE);
 
     genArray(SIZE - 1, a, 50);  // Leave a 0 at the end to test that edge case
+    //onesArray(SIZE - 1, a);
     a[SIZE - 1] = 0;
     printArray(SIZE, a, true);
 
@@ -51,48 +53,48 @@ 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
-    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, c);
+    ////printDesc("1s array for finding bugs");
+    ////StreamCompaction::Naive::scan(SIZE, c, a);
+    ////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 +139,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

@@ -23,7 +23,16 @@ namespace StreamCompaction {
          * 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
+            int tx = threadIdx.x;
+            if (tx >= n) {
+                return;
+            }
+            if (idata[tx] > 0) {
+                bools[tx] = 1;
+            }
+            else {
+                bools[tx] = 0;
+            }
         }
 
         /**
@@ -32,8 +41,13 @@ namespace StreamCompaction {
          */
         __global__ void kernScatter(int n, int *odata,
                 const int *idata, const int *bools, const int *indices) {
-            // TODO
+            int tx = threadIdx.x;
+            if (tx >= n) {
+                return;
+            }
+            if (bools[tx] == 1) {
+                odata[indices[tx]] = idata[tx];
+            }
         }
-
     }
 }

stream_compaction/cpu.cu

@@ -18,9 +18,37 @@ namespace StreamCompaction {
          * (Optional) For better understanding before starting moving to GPU, you can simulate your GPU scan in this function first.
          */
         void scan(int n, int *odata, const int *idata) {
+
+            //for (int ind = 0; ind < n / 2; ind++) {
+            //    printf("%d ", idata[ind]);
+            //}
+            //printf("\n");
+            //printf("\n");
+            //for (int ind = n / 2; ind < n; ind++) {
+            //    printf("%d ", idata[ind]);
+            //}
+            //printf("\n");
+            //printf("\n");
+
+
             timer().startCpuTimer();
-            // TODO
+            odata[0] = 0;
+            for (int ind = 1; ind < n; ind++) {
+                odata[ind] = idata[ind-1] + odata[ind - 1];
+            }
+
             timer().endCpuTimer();
+
+            //for (int ind = 0; ind < n / 2; ind++) {
+            //    printf("%d ", odata[ind]);
+            //}
+            //printf("\n");
+            //printf("\n");
+            //for (int ind = n / 2; ind < n; ind++) {
+            //    printf("%d ", odata[ind]);
+            //}
+            //printf("\n");
+            //printf("\n");
         }
 
         /**
@@ -30,9 +58,17 @@ namespace StreamCompaction {
          */
         int compactWithoutScan(int n, int *odata, const int *idata) {
             timer().startCpuTimer();
-            // TODO
+            int count = 0;
+            for (int ind = 0; ind < n; ind++) {
+                if (idata[ind] != 0) {
+                    odata[count] = idata[ind];
+                    count++;
+                }
+            }
             timer().endCpuTimer();
-            return -1;
+            return count;
+
+
         }
 
         /**
@@ -42,9 +78,36 @@ namespace StreamCompaction {
          */
         int compactWithScan(int n, int *odata, const int *idata) {
             timer().startCpuTimer();
-            // TODO
-            timer().endCpuTimer();
-            return -1;
+
+            int* bool_list = new int[n];
+            int* scan_list = new int[n];
+            for (int ind = 0; ind < n; ind++) {
+                if (idata[ind] == 0) {
+                    bool_list[ind] = 0;
+                }
+                else {
+                    bool_list[ind] = 1;
+                }
+                //printf("%d ", bool_list[ind]);
+            }
+            //printf("\n");
+            scan_list[0] = 0;
+            //printf("%d ", scan_list[0]);
+            for (int ind = 1; ind < n; ind++) {
+                scan_list[ind] = bool_list[ind - 1] + scan_list[ind - 1];
+                //printf("%d ", scan_list[ind]);
+            }
+            //printf("\n");
+            int count = 0;
+            for (int ind = 0; ind < n; ind++) {
+                //printf("%d ", idata[ind]);
+                if (bool_list[ind] == 1) {
+                    odata[scan_list[ind]] = idata[ind];
+                    count++;
+                }
+            }
+
+            return count;
         }
     }
 }