Project 2: Zwe Tun

README.md

@@ -1,14 +1,85 @@
-CUDA Stream Compaction
-======================
+# University of Pennsylvania, CIS 5650: GPU Programming and Architecture
+## Project 2 - Stream Compaction 
 
-**University of Pennsylvania, CIS 565: GPU Programming and Architecture, Project 2**
+* Zwe Tun
+  * LinkedIn: https://www.linkedin.com/in/zwe-tun-6b7191256/
+* Tested on: Intel(R) i7-14700HX, 2100 Mhz, RTX 5060 Laptop
 
-* (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)
+## Overview 
+Stream compaction is a widely used algorithm with applications in areas such as image compression, data filtering, and parallel computing. This project explores four different implementations of stream compaction:
 
-### (TODO: Your README)
+- CPU compaction without scan
 
-Include analysis, etc. (Remember, this is public, so don't put
-anything here that you don't want to share with the world.)
+- CPU compaction with scan
 
+- GPU-based compaction
+
+Each method demonstrates a different approach to optimizing performance and leveraging hardware capabilities. The goal is to analyze the performance trade-offs and efficiency of each implementation.
+
+## Implementation 
+
+###  CPU Compaction Without Scan 
+This implementation uses an iterative approach that incrementally places non-zero elements from the input array into an output array. An index counter tracks the next available position in the output. 
+
+### CPU Compaction With Scan 
+In this approach, an exclusive scan is introduced to calculate the output indices of the non-zero elements. The scan is performed using an iterative loop based on the following formula:
+
+x[i] = x[i - 1] + input[i - 1]
+
+Once the scan array is generated, we pass over the input. For each element i, if input[i] is non-zero, we place it into the output array at the position defined by scan[i].
+
+### GPU-based Compaction
+The GPU version builds on the scan-based CPU approach. Its performance depends on two scan implementations:
+
+**Naive Scan:**
+Uses GPU threads to compute the scan in multiple layers. At each layer, threads read from specific indices and write to new indices, all in place, gradually building up the scan. 
+
+**Work-Efficient Scan:**
+Uses two phases:
+
+Up-sweep: 
+Performs a parallel reduction to build a balanced binary tree of partial sums.
+
+Down-sweep:
+Traverses back down the tree to compute the exclusive scan in-place. At each pass, a node passes its value to its left child, and sets the right child to the sum of the previous left child’s value and its value. 
+
+Once the scan is complete, the result is copied back to the host (CPU), where the compaction is performed
+
+## Performance Analysis
+To evaluate and compare each implementation, a benchmark is used based on the Thrust API, specifically leveraging thrust::exclusive_scan. Performance is measured for both CPU and GPU executions by wrapping each in a performance timer class. The CPU execution time is recorded using std::chrono for high-precision timing, while GPU timing is measured using CUDA events.
+
+Note, memory allocation and transfer operations (cudaMalloc, cudaMemcpy, etc.) are excluded from the timing results.
+
+
+ 
+
+
+![Stream Compaction](img/block1.png)
+
+*Performance across various GPU block sizes for 256 array size, 128 block size used for futher tests.*
+
+
+ 
+
+
+![Stream Compaction](img/Scan.png)
+![Stream Compaction](img/compact.png)
+
+*Performance results for 8,388,608 (2^23) sized array, 128 block size.*
+
+ 
+
+
+![Stream Compaction](img/size.png)
+
+*Performance results for 8,388,608 (2^23) non-power-two vs power-of-two sized array, 128 block size.*
+
+ 
+
+## Questions 
+### To guess at what might be happening inside the Thrust implementation (e.g. allocation, memory copy), take a look at the Nsight timeline for its execution. Your analysis here doesn't have to be detailed, since you aren't even looking at the code for the implementation. Write a brief explanation of the phenomena you see here.
+
+Looking at the timeline, we can observe numerous memory operations and multiple kernel launches. These kernel launches appear to be synchronized and follow a pattern similar to the scan approaches used in this project's algorithm.
+
+### Can you find the performance bottlenecks? Is it memory I/O? Computation? Is it different for each implementation?
+For the CPU implementation, the bottleneck is predominantly computation based on the size of the array. Since the CPU executes sequentially, the main cost is due to math operations at each interation. For the GPU, memory I/O will make up the majority of the cost. Computation is parallelized through threads so global memory accesses and copies will be limiting. 

src/main.cpp

@@ -15,9 +15,9 @@
 
 const int SIZE = 1 << 8; // 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* a = new int[SIZE];
+int* b = new int[SIZE];
+int* c = new int[SIZE];
 
 int main(int argc, char* argv[]) {
     // Scan tests
@@ -29,7 +29,7 @@ int main(int argc, char* argv[]) {
 
     genArray(SIZE - 1, a, 50);  // Leave a 0 at the end to test that edge case
     a[SIZE - 1] = 0;
-    printArray(SIZE, a, true);
+    // printArray(SIZE, a, true);
 
     // initialize b using StreamCompaction::CPU::scan you implement
     // We use b for further comparison. Make sure your StreamCompaction::CPU::scan is correct.
@@ -44,7 +44,7 @@ int main(int argc, char* argv[]) {
     printDesc("cpu scan, non-power-of-two");
     StreamCompaction::CPU::scan(NPOT, c, a);
     printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)");
-    printArray(NPOT, c, true);
+    //printArray(NPOT, c, true);
     printCmpResult(NPOT, b, c);
 
     zeroArray(SIZE, c);
@@ -81,18 +81,27 @@ int main(int argc, char* argv[]) {
     //printArray(NPOT, c, true);
     printCmpResult(NPOT, b, c);
 
+   /* zeroArray(SIZE, c);
+    int* myTest = new int[SIZE] {0, 1, 2, 3, 4, 5, 6, 7};
+	int* myResult = new int[SIZE] {0, 1, 2, 6, 4, 9, 6, 28};
+    printDesc("work-efficient scan, upSweep test");
+    StreamCompaction::Efficient::scan(SIZE, c, myTest);
+    printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)");
+    printArray(SIZE, c, true);
+    printCmpResult(SIZE, myResult, 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");
@@ -104,7 +113,7 @@ int main(int argc, char* argv[]) {
 
     genArray(SIZE - 1, a, 4);  // Leave a 0 at the end to test that edge case
     a[SIZE - 1] = 0;
-    printArray(SIZE, a, true);
+    // printArray(SIZE, a, true);
 
     int count, expectedCount, expectedNPOT;
 
@@ -115,29 +124,29 @@ int main(int argc, char* argv[]) {
     count = StreamCompaction::CPU::compactWithoutScan(SIZE, b, a);
     printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(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)");
     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);
+    // 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)");
-    //printArray(count, c, true);
+    // printArray(count, c, true);
     printCmpLenResult(count, expectedCount, b, c);
 
     zeroArray(SIZE, c);
@@ -151,4 +160,4 @@ int main(int argc, char* argv[]) {
     delete[] a;
     delete[] b;
     delete[] c;
-}
+}

stream_compaction/common.cu

@@ -1,4 +1,5 @@
 #include "common.h"
+#include <device_launch_parameters.h>
 
 void checkCUDAErrorFn(const char *msg, const char *file, int line) {
     cudaError_t err = cudaGetLastError();
@@ -23,7 +24,12 @@ 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 index = threadIdx.x + (blockIdx.x * blockDim.x);
+            if (index < n) {
+                bools[index] = (idata[index] != 0) ? 1 : 0;
+			}
+
         }
 
         /**
@@ -32,7 +38,12 @@ 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 && bools[index]) {
+                odata[indices[index]] = idata[index];
+            }
+
         }
 
     }

stream_compaction/cpu.cu

@@ -19,7 +19,14 @@ 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();
         }
 
@@ -30,9 +37,17 @@ namespace StreamCompaction {
          */
         int compactWithoutScan(int n, int *odata, const int *idata) {
             timer().startCpuTimer();
-            // TODO
+
+            int count = 0; 
+            for (int i = 0; i < n; i++) {
+                if (idata[i] > 0) {
+                    odata[count] = idata[i];
+                    count++;
+                }
+            }
+
             timer().endCpuTimer();
-            return -1;
+            return count;
         }
 
         /**
@@ -42,9 +57,49 @@ namespace StreamCompaction {
          */
         int compactWithScan(int n, int *odata, const int *idata) {
             timer().startCpuTimer();
-            // TODO
+
+
+
+			int count = 0;
+			int* temp = new int[n];
+			int* scanned = new int[n];
+
+            for (int i = 0; i < n; i++) {
+                if (idata[i] != 0) {
+                    odata[i] = 1;
+                }
+                else {
+                    odata[i] = 0;
+				}
+
+
+            }
+
+
+            scanned[0] = 0;
+            for (int i = 1; i < n; i++) {
+
+                scanned[i] = scanned[i - 1] + odata[i - 1];
+            }
+
+
+
+            for (int i = 0; i < n; i++) {
+                if (odata[i] == 1) {
+
+                    temp[scanned[i]] = idata[i];
+                    count++;
+                }
+            }
+
+
+
+            for (int i = 0; i < count; i++) {
+                odata[i] = temp[i];
+            }
+
             timer().endCpuTimer();
-            return -1;
+            return count;
         }
     }
 }