Submission: SANCHIT GARG

README.md

@@ -3,211 +3,143 @@ CUDA Stream Compaction
 
 **University of Pennsylvania, CIS 565: GPU Programming and Architecture, Project 2**
 
-* (TODO) YOUR NAME HERE
-* Tested on: (TODO) Windows 22, i7-2222 @ 2.22GHz 22GB, GTX 222 222MB (Moore 2222 Lab)
+* SANCHIT GARG
+* Tested on: Mac OSX 10.10.4, i7 @ 2.4 GHz, GT 650M 1GB (Personal Computer)
 
-### (TODO: Your README)
+### SANCHIT GARG: ReadMe
 
-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 assignment, we implemented the exclusive scan and stream compaction algorithm both on the CPU and the GPU. Then we compared their performances. 
+References : http://http.developer.nvidia.com/GPUGems3/gpugems3_ch39.html
 
-Instructions (delete me)
-========================
-
-This is due Sunday, September 13 at midnight.
-
-**Summary:** In this project, you'll implement GPU stream compaction in CUDA,
-from scratch. This algorithm is widely used, and will be important for
-accelerating your path tracer project.
-
-Your stream compaction implementations in this project will simply remove `0`s
-from an array of `int`s. In the path tracer, you will remove terminated paths
-from an array of rays.
-
-In addition to being useful for your path tracer, this project is meant to
-reorient your algorithmic thinking to the way of the GPU. On GPUs, many
-algorithms can benefit from massive parallelism and, in particular, data
-parallelism: executing the same code many times simultaneously with different
-data.
-
-You'll implement a few different versions of the *Scan* (*Prefix Sum*)
-algorithm. First, you'll implement a CPU version of the algorithm to reinforce
-your understanding. Then, you'll write a few GPU implementations: "naive" and
-"work-efficient." Finally, you'll use some of these to implement GPU stream
-compaction.
+## Part 1: CPU Scan & Stream Compaction
 
-**Algorithm overview & details:** There are two primary references for details
-on the implementation of scan and stream compaction.
+A serial CPU exclusive scan and stream compaction was implemented.
 
-* The [slides on Parallel Algorithms](https://github.com/CIS565-Fall-2015/cis565-fall-2015.github.io/raw/master/lectures/2-Parallel-Algorithms.pptx)
-  for Scan, Stream Compaction, and Work-Efficient Parallel Scan.
-* GPU Gems 3, Chapter 39 - [Parallel Prefix Sum (Scan) with CUDA](http://http.developer.nvidia.com/GPUGems3/gpugems3_ch39.html).
 
-Your GPU stream compaction implementation will live inside of the
-`stream_compaction` subproject. This way, you will be able to easily copy it
-over for use in your GPU path tracer.
+## Part 2: Naive GPU Scan Algorithm
 
+A Naive GPU exclusive scan and Stream Compaction was implemented.
 
-## Part 0: The Usual
 
-This project (and all other CUDA projects in this course) requires an NVIDIA
-graphics card with CUDA capability. Any card with Compute Capability 2.0
-(`sm_20`) or greater will work. Check your GPU on this
-[compatibility table](https://developer.nvidia.com/cuda-gpus).
-If you do not have a personal machine with these specs, you may use those
-computers in the Moore 100B/C which have supported GPUs.
+## Part 3: Work-Efficient GPU Scan & Stream Compaction
 
-**HOWEVER**: If you need to use the lab computer for your development, you will
-not presently be able to do GPU performance profiling. This will be very
-important for debugging performance bottlenecks in your program.
+The Work-Efficient GPU exclusive scan and Stream Compaction was implemented.
 
-### Useful existing code
 
-* `stream_compaction/common.h`
-  * `checkCUDAError` macro: checks for CUDA errors and exits if there were any.
-  * `ilog2ceil(x)`: computes the ceiling of log2(x), as an integer.
-* `main.cpp`
-  * Some testing code for your implementations.
+## Part 4: Using Thrust's Implementation
 
+The Thrust library's exclusive scan was also implemented to compare the result with our implementations.
 
-## Part 1: CPU Scan & Stream Compaction
 
-This stream compaction method will remove `0`s from an array of `int`s.
+## Part 5: Radix Sort
 
-In `stream_compaction/cpu.cu`, implement:
+Implemented the Parallel Radix Sort algorithm as explained in the reference.
+A Namespace "Radix"
 
-* `StreamCompaction::CPU::scan`: compute an exclusive prefix sum.
-* `StreamCompaction::CPU::compactWithoutScan`: stream compaction without using
-  the `scan` function.
-* `StreamCompaction::CPU::compactWithScan`: stream compaction using the `scan`
-  function. Map the input array to an array of 0s and 1s, scan it, and use
-  scatter to produce the output. You will need a **CPU** scatter implementation
-  for this (see slides or GPU Gems chapter for an explanation).
+## Performance Analysis
 
-These implementations should only be a few lines long.
+My implementation observed the following pattern. The time are all in milliseconds. I used 1024 threads per block for all GPU implementation
 
+#### Values
 
-## Part 2: Naive GPU Scan Algorithm
+![](images/Values.png)
 
-In `stream_compaction/naive.cu`, implement `StreamCompaction::Naive::scan`
+#### Graph
 
-This uses the "Naive" algorithm from GPU Gems 3, Section 39.2.1. We haven't yet
-taught shared memory, and you **shouldn't use it yet**. Example 39-1 uses
-shared memory, but is limited to operating on very small arrays! Instead, write
-this using global memory only. As a result of this, you will have to do
-`ilog2ceil(n)` separate kernel invocations.
+![](images/PerformanceGraph.png)
 
-Beware of errors in Example 39-1 in the book; both the pseudocode and the CUDA
-code in the online version of Chapter 39 are known to have a few small errors
-(in superscripting, missing braces, bad indentation, etc.)
+#### Analysis
 
-Since the parallel scan algorithm operates on a binary tree structure, it works
-best with arrays with power-of-two length. Make sure your implementation works
-on non-power-of-two sized arrays (see `ilog2ceil`). This requires extra memory
-- your intermediate array sizes will need to be rounded to the next power of
-two.
+The bottleneck in the Naive implementation would be copying the output array after every scan step to use it as the input array. Switching the arrays was giving incorrect results.
+The bottleneck for the Work-Efficient implementation should be a lot of memory access in the kernel functions. This is a slow process and hence reduces the performance of the implementation. 
 
+### Output
 
-## Part 3: Work-Efficient GPU Scan & Stream Compaction
+The console output of the program is as follows. Note that a test for Radix Sort was also written. I used the sort function under the header file "algorithm" and compared my implementations result with it. This was done to test the correctness of the Parallel Radix Sort implementation.
 
-### 3.1. Scan
+****************
+** SCAN TESTS **
+****************
 
-In `stream_compaction/efficient.cu`, implement
-`StreamCompaction::Efficient::scan`
+    [  30  41  15  22  11  41  10  37  48  41  44  30  26 ...  20   0 ]
+==== cpu scan, power-of-two ====
 
-All of the text in Part 2 applies.
+    [   0  30  71  86 108 119 160 170 207 255 296 340 370 ... 26119 26139 ]
+==== cpu scan, non-power-of-two ====
 
-* This uses the "Work-Efficient" algorithm from GPU Gems 3, Section 39.2.2.
-* Beware of errors in Example 39-2.
-* Test non-power-of-two sized arrays.
+    [   0  30  71  86 108 119 160 170 207 255 296 340 370 ... 26031 26064 ]
+    passed 
 
-### 3.2. Stream Compaction
+==== naive scan, power-of-two ====
 
-This stream compaction method will remove `0`s from an array of `int`s.
+    [   0  30  71  86 108 119 160 170 207 255 296 340 370 ... 26119 26139 ]
+    passed 
 
-In `stream_compaction/efficient.cu`, implement
-`StreamCompaction::Efficient::compact`
+==== naive scan, non-power-of-two ====
 
-For compaction, you will also need to implement the scatter algorithm presented
-in the slides and the GPU Gems chapter.
+    [   0  30  71  86 108 119 160 170 207 255 296 340 370 ...   0   0 ]
+    passed 
 
-In `stream_compaction/common.cu`, implement these for use in `compact`:
+==== work-efficient scan, power-of-two ====
 
-* `StreamCompaction::Common::kernMapToBoolean`
-* `StreamCompaction::Common::kernScatter`
+    [   0  30  71  86 108 119 160 170 207 255 296 340 370 ... 26119 26139 ]
+    passed 
 
+==== work-efficient scan, non-power-of-two ====
 
-## Part 4: Using Thrust's Implementation
+    [   0  30  71  86 108 119 160 170 207 255 296 340 370 ... 26031 26064 ]
+    passed 
 
-In `stream_compaction/thrust.cu`, implement:
+==== thrust scan, power-of-two ====
 
-* `StreamCompaction::Thrust::scan`
+    [   0  30  71  86 108 119 160 170 207 255 296 340 370 ... 26119 26139 ]
+    passed 
 
-This should be a very short function which wraps a call to the Thrust library
-function `thrust::exclusive_scan(first, last, result)`.
+==== thrust scan, non-power-of-two ====
 
-To measure timing, be sure to exclude memory operations by passing
-`exclusive_scan` a `thrust::device_vector` (which is already allocated on the
-GPU).  You can create a `thrust::device_vector` by creating a
-`thrust::host_vector` from the given pointer, then casting it.
+    [   0  30  71  86 108 119 160 170 207 255 296 340 370 ... 26031 26064 ]
+    passed 
 
 
-## Part 5: Radix Sort (Extra Credit) (+10)
+*****************************
+** STREAM COMPACTION TESTS **
+*****************************
 
-Add an additional module to the `stream_compaction` subproject. Implement radix
-sort using one of your scan implementations. Add tests to check its correctness.
+    [   2   3   3   0   1   1   2   1   2   1   2   0   2 ...   0   0 ]
 
+==== cpu compact without scan, power-of-two ====
 
-## Write-up
+    [   2   3   3   1   1   2   1   2   1   2   2   2   3 ...   2   1 ]
+    passed 
 
-1. Update all of the TODOs at the top of this README.
-2. Add a description of this project including a list of its features.
-3. Add your performance analysis (see below).
+==== cpu compact without scan, non-power-of-two ====
 
-All extra credit features must be documented in your README, explaining its
-value (with performance comparison, if applicable!) and showing an example how
-it works. For radix sort, show how it is called and an example of its output.
+    [   2   3   3   1   1   2   1   2   1   2   2   2   3 ...   1   2 ]
+    passed 
 
-Always profile with Release mode builds and run without debugging.
+==== cpu compact with scan ====
 
-### Questions
+    [   2   3   3   1   1   2   1   2   1   2   2   2   3 ...   2   1 ]
+    passed 
 
-* Roughly optimize the block sizes of each of your implementations for minimal
-  run time on your GPU.
-  * (You shouldn't compare unoptimized implementations to each other!)
+==== work-efficient compact, power-of-two ====
 
-* Compare all of these GPU Scan implementations (Naive, Work-Efficient, and
-  Thrust) to the serial CPU version of Scan. Plot a graph of the comparison
-  (with array size on the independent axis).
-  * You should use CUDA events for timing. Be sure **not** to include any
-    explicit memory operations in your performance measurements, for
-    comparability.
-  * To guess at what might be happening inside the Thrust implementation, take
-    a look at the Nsight timeline for its execution.
+    [   2   3   3   1   1   2   1   2   1   2   2   2   3 ...   2   1 ]
+    passed 
 
-* Write a brief explanation of the phenomena you see here.
-  * Can you find the performance bottlenecks? Is it memory I/O? Computation? Is
-    it different for each implementation?
+==== work-efficient compact, non-power-of-two ====
 
-* Paste the output of the test program into a triple-backtick block in your
-  README.
-  * If you add your own tests (e.g. for radix sort or to test additional corner
-    cases), be sure to mention it explicitly.
+    [   2   3   3   1   1   2   1   2   1   2   2   2   3 ...   1   2 ]
+    passed 
 
-These questions should help guide you in performance analysis on future
-assignments, as well.
+****************
+** RADIX SORT **
+****************
 
-## Submit
+==== Radix Sort, sizeAr elements ====
 
-If you have modified any of the `CMakeLists.txt` files at all (aside from the
-list of `SOURCE_FILES`), you must test that your project can build in Moore
-100B/C. Beware of any build issues discussed on the Google Group.
+    [  30  91  15  72  61  41  10  37  98  41  94  80  26  96  10  88 ]
+
+    [  10  10  15  26  30  37  41  41  61  72  80  88  91  94  96  98 ]
+    passed 
 
-1. Open a GitHub pull request so that we can see that you have finished.
-   The title should be "Submission: YOUR NAME".
-2. Send an email to the TA (gmail: kainino1+cis565@) with:
-   * **Subject**: in the form of `[CIS565] Project 2: PENNKEY`
-   * Direct link to your pull request on GitHub
-   * In the form of a grade (0-100+) with comments, evaluate your own
-     performance on the project.
-   * Feedback on the project itself, if any.

cis565_stream_compaction_test.launch

@@ -8,8 +8,8 @@
 <stringAttribute key="org.eclipse.cdt.launch.DEBUGGER_START_MODE" value="run"/>
 <stringAttribute key="org.eclipse.cdt.launch.PROGRAM_NAME" value="build/cis565_stream_compaction_test"/>
 <stringAttribute key="org.eclipse.cdt.launch.PROJECT_ATTR" value="Project2-Stream-Compaction"/>
-<booleanAttribute key="org.eclipse.cdt.launch.PROJECT_BUILD_CONFIG_AUTO_ATTR" value="true"/>
-<stringAttribute key="org.eclipse.cdt.launch.PROJECT_BUILD_CONFIG_ID_ATTR" value=""/>
+<booleanAttribute key="org.eclipse.cdt.launch.PROJECT_BUILD_CONFIG_AUTO_ATTR" value="false"/>
+<stringAttribute key="org.eclipse.cdt.launch.PROJECT_BUILD_CONFIG_ID_ATTR" value="com.nvidia.cuda.ide.toolchain.base.1399573849.1760580076"/>
 <booleanAttribute key="org.eclipse.cdt.launch.use_terminal" value="true"/>
 <listAttribute key="org.eclipse.debug.core.MAPPED_RESOURCE_PATHS">
 <listEntry value="/Project2-Stream-Compaction"/>
@@ -18,8 +18,8 @@
 <listEntry value="4"/>
 </listAttribute>
 <listAttribute key="org.eclipse.debug.ui.favoriteGroups">
-<listEntry value="org.eclipse.debug.ui.launchGroup.profile"/>
 <listEntry value="org.eclipse.debug.ui.launchGroup.debug"/>
+<listEntry value="org.eclipse.debug.ui.launchGroup.profile"/>
 <listEntry value="org.eclipse.debug.ui.launchGroup.run"/>
 </listAttribute>
 <stringAttribute key="org.eclipse.dsf.launch.MEMORY_BLOCKS" value="&lt;?xml version=&quot;1.0&quot; encoding=&quot;UTF-8&quot; standalone=&quot;no&quot;?&gt;&#10;&lt;memoryBlockExpressionList context=&quot;reserved-for-future-use&quot;/&gt;&#10;"/>