CIS565-Fall-2016 · itoupeter · Oct 17, 2016 · Oct 22, 2016 · Oct 22, 2016 · Oct 23, 2016
diff --git a/README.md b/README.md
@@ -5,14 +5,78 @@ CUDA Rasterizer
 
 **University of Pennsylvania, CIS 565: GPU Programming and Architecture, Project 4**
 
-* (TODO) YOUR NAME HERE
-* Tested on: (TODO) Windows 22, i7-2222 @ 2.22GHz 22GB, GTX 222 222MB (Moore 2222 Lab)
+* Liang Peng
+* Tested on: Windows 10, i7-6700HQ @ 2.6GHz 2.6GHz 8GB, GTX 960M (Personal Laptop)
 
-### (TODO: Your README)
+## Features
+* [x] Primitive
+	* [x] Point
+	* [x] Line
+	* [x] Triangle
+* [x] Texture
+	* [x] Diffuse
+	* [x] Bilinear Filtering
+* [x] Lighting
+	* [x] Lambert
+	* [x] Blinn-Phong
+* [x] Normal Visualization
+* [x] Depth Visualization
+* [x] Depth Test with Mutex
+* [x] Perspective Corrected Texcoord
+* [x] Performance Analysis
 
-*DO NOT* leave the README to the last minute! It is a crucial part of the
-project, and we will not be able to grade you without a good README.
+## Overview
+### Basic Attributes
+Texture | Normal | Depth | Texcoord
+--- | --- | --- | ---
+![](img/cover_diffuse.gif) | ![](img/cover_normal.gif) | ![](img/cover_depth.gif) | ![](img/cover_texcoord.gif)
+_Note*_ Attributes displayed on each pixel are calculated by several steps. Vertex Shader processes each vertex by transforming vertex positions from model space to view space then to projection space, calculating eye positions, eye normals, texcoords. Next Primitive assembler will group all vertices in primitives, in this project, triangles. Then rasterizer will proccess every triangle by testing every pixel in its bounding box. If a pixel is on the triangle, its position, normal, texcoord are interpolated using barycentric weights on the threes vertices, and stored in fragment buffer. Last Fragment Shader will calculate color to shade each pixel using these attributes and store the color in frame buffer.
 
+### Lighting
+Lambert | Blinn-Phong
+--- | ---
+ ![](img/cover_lambert.gif) | ![](img/cover_blinnphong.gif)
+_Note*_ Lighting is calculated in Fragment Shader for each pixel. Lambert reflectance is proportional to the dot product of surface normal and light direction, while Blinn-Phong reflectance is proportional to the dot product of surface normal and halfway direction raised to the power of shininess.
+
+### Resterization Mode
+Point Cloud | Wireframe | Solid
+--- | --- | ---
+![](img/point.gif) | ![](img/line.gif) | ![](img/triangle.gif)
+ _Note*_ Point Cloud mode is implemented in a straight forward way, by shading the pixel where each pixel is located. Wireframe mode is most concerned with how to shade the pixels between two vertices. In this project, I implemented that by enforcing one pixel in each row if a segment is more vertical than 45 degree, or one pixel in each column if the segment is more horizontal than 45 degree.
+
+### Depth Test with Mutex
+Mutex OFF | Mutex ON
+:---:|:---:
+63 ms/frame | 74 ms/frame
+![](img/mutex_off.gif) | ![](img/mutex_on.gif)
+ _Note*_ We can see in the animated GIFs, noise appears in Mutex OFF mode, because of race condition. To avoid that, each pixel has its own mutex, which is essentially an unsigned integer. When a CUDA thread is trying to seize a pixel to write, it calls atomicCAS to try to change the value of mutex. If it succeeded, it can then safely write that pixel, or it would try again. It alse takes more time to process a frame with Mutex.
+
+### Texcoord Correction
+Perspective Correction OFF | Perspective Correction ON
+:---:|:---:
+ ![](img/texcoord0.gif) | ![](img/texcoord1.gif)
+ _Note*_ Texcoord Correction is unconditionally necessary since texture mapping would look wrong without correction. This is due to the fact that texture coordinates are interpolated in screen space, ignoring the depth information (Texcoords with deeper depth would be denser).
+
+### Texture Filtering
+Nearest Neighbor | Bilinear
+--- | ---
+ ![](img/nearestneighbor.gif) | ![](img/bilinear.gif)
+ _Note*_ Bilinear filtering is useful because texcoords are normalized thus cannot be mapped to a texel exactly after being scaled by texture width and height, thus we grab the four texels around the scaled texcoord and bilinear interpolate the color.
+
+### Performance Analysis
+* Resterization Mode
+
+![](img/chart1.png)
+
+_Note*_ From the stacked column chart we can see that point cloud mode has the lowest percentage of rasterization runtime, while solid mode has the highest percentage. This is because in solid mode, every pixel within the bounding box of each triangle is tested, which makes rasterization of triangles computation-consuming.
+
+* Different Models
+
+Cow | Duck | Truck
+--- | --- | ---
+<img src="img/cow.PNG" height="200"> | <img src="img/duck.PNG" height="200"> | <img src="img/truck.PNG" height="200">
+
+![](img/chart2.png)
 
 ### Credits
 

diff --git a/img/bilinear.gif b/img/bilinear.gif
diff --git a/img/chart1.png b/img/chart1.png
diff --git a/img/chart2.png b/img/chart2.png
diff --git a/img/cover_blinnphong.gif b/img/cover_blinnphong.gif
diff --git a/img/cover_depth.gif b/img/cover_depth.gif
diff --git a/img/cover_diffuse.gif b/img/cover_diffuse.gif
diff --git a/img/cover_lambert.gif b/img/cover_lambert.gif
diff --git a/img/cover_normal.gif b/img/cover_normal.gif
diff --git a/img/cover_texcoord.gif b/img/cover_texcoord.gif
diff --git a/img/cow.PNG b/img/cow.PNG
diff --git a/img/duck.PNG b/img/duck.PNG
diff --git a/img/dummy.gif b/img/dummy.gif
diff --git a/img/line.gif b/img/line.gif
diff --git a/img/mutex_off.gif b/img/mutex_off.gif
diff --git a/img/mutex_on.gif b/img/mutex_on.gif
diff --git a/img/nearestneighbor.gif b/img/nearestneighbor.gif
diff --git a/img/point.gif b/img/point.gif
diff --git a/img/texcoord0.gif b/img/texcoord0.gif
diff --git a/img/texcoord1.gif b/img/texcoord1.gif
diff --git a/img/triangle.gif b/img/triangle.gif
diff --git a/img/truck.PNG b/img/truck.PNG
diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt
@@ -6,5 +6,5 @@ set(SOURCE_FILES
 
 cuda_add_library(src
     ${SOURCE_FILES}
-    OPTIONS -arch=sm_20
+    OPTIONS -arch=sm_50
     )
diff --git a/src/main.cpp b/src/main.cpp
@@ -13,6 +13,7 @@
 #define STB_IMAGE_IMPLEMENTATION
 #define TINYGLTF_LOADER_IMPLEMENTATION
 #include <util/tiny_gltf_loader.h>
+#include <chrono>
 
 //-------------------------------
 //-------------MAIN--------------
@@ -65,18 +66,26 @@ int main(int argc, char **argv) {
 void mainLoop() {
     while (!glfwWindowShouldClose(window)) {
         glfwPollEvents();
+		// tic
+		auto tic = std::chrono::high_resolution_clock::now();
         runCuda();
+		cudaDeviceSynchronize();
+		// toc
+		auto toc = std::chrono::high_resolution_clock::now();
+		std::chrono::milliseconds tictoc =
+				std::chrono::duration_cast<std::chrono::milliseconds>(toc - tic);
 
         time_t seconds2 = time (NULL);
 
         if (seconds2 - seconds >= 1) {
-
             fps = fpstracker / (seconds2 - seconds);
             fpstracker = 0;
             seconds = seconds2;
         }
 
-        string title = "CIS565 Rasterizer | " + utilityCore::convertIntToString((int)fps) + " FPS";
+        string title = "CIS565 Rasterizer | "
+				+ utilityCore::convertIntToString((int)fps) + " FPS | ["
+				+ utilityCore::convertIntToString(tictoc.count()) + " ms]";
         glfwSetWindowTitle(window, title.c_str());
 
         glBindBuffer(GL_PIXEL_UNPACK_BUFFER, pbo);
@@ -95,7 +104,7 @@ void mainLoop() {
 //-------------------------------
 //---------RUNTIME STUFF---------
 //-------------------------------
-float scale = 1.0f;
+float scale = 1.f;
 float x_trans = 0.0f, y_trans = 0.0f, z_trans = -10.0f;
 float x_angle = 0.0f, y_angle = 0.0f;
 void runCuda() {

diff --git a/src/main.hpp b/src/main.hpp
@@ -101,4 +101,4 @@ std::string getFilePathExtension(const std::string &FileName);
 
 void mouseButtonCallback(GLFWwindow* window, int button, int action, int mods);
 void mouseMotionCallback(GLFWwindow* window, double xpos, double ypos);
-void mouseWheelCallback(GLFWwindow* window, double xoffset, double yoffset);
+void mouseWheelCallback(GLFWwindow* window, double xoffset, double yoffset);