Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
84 changes: 78 additions & 6 deletions README.md
Original file line number Diff line number Diff line change
Expand Up @@ -3,12 +3,84 @@ 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)
* Hanyu Liu
* [personal website](http://liuhanyu.net/)
* Tested on: Windows 10, Intel(R) Core(TM) i7-7700HQ CPU @ 2.80GHz 16.0GB, GeForce GTX 1050 (Personal)

### Summary of Project

In this project, I implemented stream compaction, which simply removes `0`s from an array of `1`s, on both the CPU and the GPU in CUDA. Stream compaction uses the scan (Prefix Sum) Algorithm, and in this project, I implemented four different scan methods: 1) Scan on the CPU, 2) Naive, 3) Work-Efficient, and 4) Using Thrust. Furthermore, GPU stream compaction also needed additional kernels to generate a boolean mask and scatter, which I also implemented. With the help of these functions, I was able to implement stream compaction on both the GPU and the CPU.



Stream compaction is widely used, and will be used to accelerate my future path tracer project.



### Performance Analysis

![](img/performance.png)

1. For the Thrust implementation, the runtime is significantly lower than the other implementations. This is possibly due to the small amount of memory copy as it alters the data in place without the need for extra buffers.

2. Here, we see that the Naive and Work-Efficient scans both take much longer than the CPU scan at large array sizes even though we are altering the array in parallel on the GPU. In fact, the work-efficient scan takes much longer than the naive scan. Both GPU implementations take longer possibly because there is more overhead from the kernel calls. As for the difference between Naive and Work-Efficient scans, the work-efficient scan takes two for loops on larger, padded arrays, which causes it to be slower than the naive implementation. Ultimately, the bottle-neck is the number of additions we have to perform, which is log base2 of n. The rest of the performance depends on memory allocation, the cache, and overhead, in which case, CPU would win out.

3. ```

****************
** SCAN TESTS **
****************
[ 3 5 18 11 7 47 39 47 30 45 5 9 40 ... 23 0 ]
==== cpu scan, power-of-two ====
elapsed time: 0.0023ms (std::chrono Measured)
[ 0 3 8 26 37 44 91 130 177 207 252 257 266 ... 24162 24185 ]
==== cpu scan, non-power-of-two ====
elapsed time: 0.0051ms (std::chrono Measured)
[ 0 3 8 26 37 44 91 130 177 207 252 257 266 ... 24142 24149 ]
passed
==== naive scan, power-of-two ====
elapsed time: 0.023552ms (CUDA Measured)
passed
==== naive scan, non-power-of-two ====
elapsed time: 0.022528ms (CUDA Measured)
passed
==== work-efficient scan, power-of-two ====
elapsed time: 0.1024ms (CUDA Measured)
passed
==== work-efficient scan, non-power-of-two ====
elapsed time: 0.132096ms (CUDA Measured)
passed
==== thrust scan, power-of-two ====
elapsed time: 0.079872ms (CUDA Measured)
passed
==== thrust scan, non-power-of-two ====
elapsed time: 0.090336ms (CUDA Measured)
passed

*****************************
** STREAM COMPACTION TESTS **
*****************************
[ 3 3 2 1 3 3 3 3 0 1 3 3 2 ... 1 0 ]
==== cpu compact without scan, power-of-two ====
elapsed time: 0.0073ms (std::chrono Measured)
[ 3 3 2 1 3 3 3 3 1 3 3 2 2 ... 2 1 ]
passed
==== cpu compact without scan, non-power-of-two ====
elapsed time: 0.0074ms (std::chrono Measured)
[ 3 3 2 1 3 3 3 3 1 3 3 2 2 ... 3 1 ]
passed
==== cpu compact with scan ====
elapsed time: 0.0363ms (std::chrono Measured)
[ 3 3 2 1 3 3 3 3 1 3 3 2 2 ... 2 1 ]
passed
==== work-efficient compact, power-of-two ====
elapsed time: 0.192512ms (CUDA Measured)
passed
==== work-efficient compact, non-power-of-two ====
elapsed time: 0.171008ms (CUDA Measured)
passed
Press any key to continue . . .
```

### (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.)

Binary file added img/performance.png
Loading
Sorry, something went wrong. Reload?
Sorry, we cannot display this file.
Sorry, this file is invalid so it cannot be displayed.
2 changes: 1 addition & 1 deletion src/main.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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 = 1000; // 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];
Expand Down
17 changes: 15 additions & 2 deletions stream_compaction/common.cu
Original file line number Diff line number Diff line change
Expand Up @@ -23,7 +23,15 @@ 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 k = (blockIdx.x * blockDim.x) + threadIdx.x;
if (k >= n) return;

if (idata[k] != 0) {
bools[k] = 1;
}
else {
bools[k] = 0;
}
}

/**
Expand All @@ -32,7 +40,12 @@ namespace StreamCompaction {
*/
__global__ void kernScatter(int n, int *odata,
const int *idata, const int *bools, const int *indices) {
// TODO
int k = (blockIdx.x * blockDim.x) + threadIdx.x;
if (k >= n) return;

if (bools[k] == 1) {
odata[indices[k]] = idata[k];
}
}

}
Expand Down
57 changes: 52 additions & 5 deletions stream_compaction/cpu.cu
Original file line number Diff line number Diff line change
Expand Up @@ -19,7 +19,11 @@ namespace StreamCompaction {
*/
void scan(int n, int *odata, const int *idata) {
timer().startCpuTimer();
// TODO
// compute an exclusive prefix sum
odata[0] = 0;
for (int i = 1; i < n; i++) {
odata[i] = odata[i - 1] + idata[i - 1];
}
timer().endCpuTimer();
}

Expand All @@ -30,9 +34,28 @@ namespace StreamCompaction {
*/
int compactWithoutScan(int n, int *odata, const int *idata) {
timer().startCpuTimer();
// TODO
int elements_remaining = 0;
for (int i = 0; i < n; i++) {
if (idata[i] != 0) {
odata[elements_remaining] = idata[i];
elements_remaining++;
}
}
timer().endCpuTimer();
return -1;
return elements_remaining;
}

/*
* Helper Function because I seem to be having issues when I start the timer again
* Same as scan function earlier, just without the timer
* From Piazza @110
*/
void cpu_scan(int n, int* odata, const int* idata) {
// compute an exclusive prefix sum
odata[0] = 0;
for (int i = 1; i < n; i++) {
odata[i] = odata[i - 1] + idata[i - 1];
}
}

/**
Expand All @@ -42,9 +65,33 @@ namespace StreamCompaction {
*/
int compactWithScan(int n, int *odata, const int *idata) {
timer().startCpuTimer();
// TODO
// fill temp array with 0 if idata is 0 or 1 otherwise
int* temp_array = new int[n];
for (int i = 0; i < n; i++) {
if (idata[i] == 0) {
temp_array[i] = 0;
}
else temp_array[i] = 1;
}

// run exclusive scan on temporary array
int* scanned_array = new int[n];
cpu_scan(n, scanned_array, temp_array);

// scatter
for (int j = 0; j < n; j++) {
if (temp_array[j] == 1) {
// write element
odata[scanned_array[j]] = idata[j];
}
}

// cleanup
int result = scanned_array[n - 1];
timer().endCpuTimer();
return -1;
delete[] temp_array;
delete[] scanned_array;
return result;
}
}
}
Loading