Project 2: Jason Xie

.vscode/launch.json

@@ -0,0 +1,20 @@
+{
+    // Use IntelliSense to learn about possible attributes.
+    // Hover to view descriptions of existing attributes.
+    // For more information, visit: https://go.microsoft.com/fwlink/?linkid=830387
+    "version": "0.2.0",
+    "configurations": [
+        {
+            "name": "CUDA C++: Launch",
+            "type": "cuda-gdb",
+            "request": "launch",
+            "program": ""
+        },
+        {
+            "name": "CUDA C++: Attach",
+            "type": "cuda-gdb",
+            "request": "attach",
+            "processId": "${command:cuda.pickProcess}"
+        }
+    ]
+}

.vscode/settings.json

@@ -0,0 +1,18 @@
+{
+    "ros.distro": "humble",
+    "files.associations": {
+        "*.sdf": "xml",
+        "*.world": "xml",
+        "array": "cpp",
+        "deque": "cpp",
+        "forward_list": "cpp",
+        "list": "cpp",
+        "string": "cpp",
+        "unordered_map": "cpp",
+        "unordered_set": "cpp",
+        "vector": "cpp",
+        "string_view": "cpp",
+        "initializer_list": "cpp",
+        "numeric": "cpp"
+    }
+}

README.md

@@ -1,14 +1,101 @@
 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)
+**Jason Xie**
 
-### (TODO: Your README)
+[🤓 LinkedIn](https://linkedin.com/in/jia-chun-xie)
 
-Include analysis, etc. (Remember, this is public, so don't put
-anything here that you don't want to share with the world.)
+[😇 my website](https://jchunx.dev)
 
+[🥵 X (formerly 🐦)](https://x.com/codemonke_)
+
+Tested on: Ubuntu 22.04, i5-8400, RTX 3060Ti, personal machine
+
+
+## About
+
+CUDA implementation of stream compaction & scan.
+
+## Performance Analysis
+
+### Block Size
+
+For all kernels, I found that 128 was the optimal block size.
+
+### Time vs. Array Size
+
+![comparison plot](assets/cuda-perf-compare-proj2.png)
+| Array Size | CPU Scan | Naive Scan | Work-Efficient Scan | Thrust Scan | CPU Compact | Work-Efficient Compact |
+| ---------- | -------- | ---------- | ------------------- | ----------- | ----------- | ---------------------- |
+| 2^8        | 0.000333 | 0.0806 | 0.0963 | 0.0522 | 0.000834 | 0.125 |
+| 2^12 | 0.00235 | 0.0840 | 0.109 | 0.0543 | 0.00924 | 0.108 |
+| 2^16 | 0.0388 | 0.0820 | 0.120 | 0.0522 | 0.145 | 0.107 |
+| 2^20 | 0.541 | 0.503 | 0.533 | 0.194 | 2.42 | 0.726 |
+| 2^24 | 9.033 | 11.134 | 7.661 | 0.500 | 38.880 | 8.487 |
+
+## What is going on here?
+
+Unoptimized GPU scanning actually performs worse than CPU scanning. A bit of Nsight shows that the kernels have low warp occupancy (esp. up / down sweeps):
+![low occupancy](assets/low-occupancy.png)
+
+The hypothesis here is that the kernels are being bottlenecked by global memory access and warp divergence.
+
+## Test Program Outputs
+
+```
+****************
+** SCAN TESTS **
+****************
+    [  38  28  32  41  11   4  49  31  48   8  42  48  22 ...  25   0 ]
+==== cpu scan, power-of-two ====
+   elapsed time: 4.48317ms    (std::chrono Measured)
+    [  38  66  98 139 150 154 203 234 282 290 332 380 402 ... 205516747 205516747 ]
+==== cpu scan, non-power-of-two ====
+   elapsed time: 4.44224ms    (std::chrono Measured)
+    [  38  66  98 139 150 154 203 234 282 290 332 380 402 ... 205516679 205516703 ]
+    passed 
+==== naive scan, power-of-two ====
+   elapsed time: 4.82586ms    (CUDA Measured)
+    a[4194304] = 102760269, b[4194304] = 102760231
+    FAIL VALUE 
+==== naive scan, non-power-of-two ====
+   elapsed time: 4.18042ms    (CUDA Measured)
+    a[4194304] = 102760269, b[4194304] = 102760231
+    FAIL VALUE 
+==== work-efficient scan, power-of-two ====
+   elapsed time: 3.68918ms    (CUDA Measured)
+    passed 
+==== work-efficient scan, non-power-of-two ====
+   elapsed time: 3.61574ms    (CUDA Measured)
+    passed 
+==== thrust scan, power-of-two ====
+   elapsed time: 0.346912ms    (CUDA Measured)
+    passed 
+==== thrust scan, non-power-of-two ====
+   elapsed time: 0.310976ms    (CUDA Measured)
+    passed 
+
+*****************************
+** STREAM COMPACTION TESTS **
+*****************************
+    [   0   0   2   1   1   0   3   1   2   0   2   2   0 ...   3   0 ]
+==== cpu compact without scan, power-of-two ====
+   elapsed time: 19.3632ms    (std::chrono Measured)
+    [   2   1   1   3   1   2   2   2   2   2   2   1   3 ...   3   3 ]
+    passed 
+==== cpu compact without scan, non-power-of-two ====
+   elapsed time: 19.8116ms    (std::chrono Measured)
+    [   2   1   1   3   1   2   2   2   2   2   2   1   3 ...   3   2 ]
+    passed 
+==== cpu compact with scan ====
+   elapsed time: 71.7739ms    (std::chrono Measured)
+    [   2   1   1   3   1   2   2   2   2   2   2   1   3 ...   3   3 ]
+    passed 
+==== work-efficient compact, power-of-two ====
+   elapsed time: 4.29862ms    (CUDA Measured)
+    passed 
+==== work-efficient compact, non-power-of-two ====
+   elapsed time: 4.53734ms    (CUDA Measured)
+    passed 
+```

build.sh

@@ -0,0 +1,26 @@
+#!/bin/bash
+
+# Create build directory if it doesn't exist
+mkdir -p build
+
+# Change into the build directory
+cd build
+
+# Default build type to Release
+build_type="Release"
+
+# Check for argument "debug" to change build type
+if [ "$1" == "debug" ]; then
+  echo "Building in debug mode"
+  build_type="Debug"
+fi
+
+# Run cmake with the specified build type
+cmake -DCMAKE_BUILD_TYPE=$build_type ..
+
+# Build the project with dbg if debug was specified
+if [ "$1" == "debug" ]; then
+  make dbg=1
+else
+  make
+fi

src/main.cpp

@@ -11,17 +11,20 @@
 #include <stream_compaction/naive.h>
 #include <stream_compaction/efficient.h>
 #include <stream_compaction/thrust.h>
+#include <unistd.h>
 #include "testing_helpers.hpp"
 
-const int SIZE = 1 << 8; // feel free to change the size of array
+const int SIZE = 1 << 24; // 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 main(int argc, char* argv[]) {
     // Scan tests
-
+    // std::cout << "sleeping..." << std::endl;
+    // sleep(5);
+    // std::cout << "awake!" << std::endl;
     printf("\n");
     printf("****************\n");
     printf("** SCAN TESTS **\n");

stream_compaction/common.cu

@@ -23,7 +23,11 @@ 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) {
+                return;
+            }
+            bools[index] = idata[index] == 0 ? 0 : 1;
         }
 
         /**
@@ -32,7 +36,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] == 1) {
+                odata[indices[index]] = idata[index];
+            } 
         }
 
     }

stream_compaction/cpu.cu

@@ -19,20 +19,40 @@ namespace StreamCompaction {
          */
         void scan(int n, int *odata, const int *idata) {
             timer().startCpuTimer();
-            // TODO
+            for (int i = 0; i < n; i++) {
+                odata[i] = idata[i];
+                if (i > 0) {
+                    odata[i] += odata[i - 1];
+                }
+            }
             timer().endCpuTimer();
         }
 
+        void _scan_no_timer(int n, int *odata, const int *idata) {
+            for (int i = 0; i < n; i++) {
+                odata[i] = idata[i];
+                if (i > 0) {
+                    odata[i] += odata[i - 1];
+                }
+            }
+        }
+
         /**
          * CPU stream compaction without using the scan function.
          *
          * @returns the number of elements remaining after compaction.
          */
         int compactWithoutScan(int n, int *odata, const int *idata) {
             timer().startCpuTimer();
-            // TODO
+            int numElements = 0;
+            for (int i = 0; i < n; i++) {
+                if (idata[i] != 0) {
+                    odata[numElements] = idata[i];
+                    numElements++;
+                }
+            }
             timer().endCpuTimer();
-            return -1;
+            return numElements;
         }
 
         /**
@@ -42,9 +62,29 @@ namespace StreamCompaction {
          */
         int compactWithScan(int n, int *odata, const int *idata) {
             timer().startCpuTimer();
-            // TODO
+            int* bools = new int[n];
+            for (int i = 0; i < n; i++) {
+                bools[i] = (idata[i] != 0) ? 1 : 0;
+            }
+            int* scanned = new int[n];
+            _scan_no_timer(n, scanned, bools);
+            // convert to exclusive scan
+            for (int i = n - 1; i > 0; i--) {
+                scanned[i] = scanned[i - 1];
+            }
+            scanned[0] = 0;
+            int numElements = 0;
+            for (int i = 0; i < n; i++) {
+                if (bools[i] != 0) {
+                    odata[scanned[i]] = idata[i];
+                    numElements++;
+                }
+            }
+
+            delete[] bools;
+            delete[] scanned;
             timer().endCpuTimer();
-            return -1;
+            return numElements;
         }
     }
 }