diff --git a/.cproject b/.cproject
index 8757ba5..a2a201f 100644
--- a/.cproject
+++ b/.cproject
@@ -21,7 +21,7 @@
 					<folderInfo id="com.nvidia.cuda.ide.toolchain.on_path.1673801701.1152097236" name="/" resourcePath="">
 						<toolChain id="com.nvidia.cuda.ide.toolchain.on_path.1386955599" name="CUDA Toolkit on the path" superClass="com.nvidia.cuda.ide.toolchain.on_path">
 							<targetPlatform archList="all" binaryParser="com.nvidia.cuda.ide.elf;com.nvidia.cuda.ide.macho;com.nvidia.cuda.ide.cubin" id="com.nvidia.cuda.ide.targetPlatform.853269530" isAbstract="false" name="Debug Platform" osList="linux,macosx" superClass="com.nvidia.cuda.ide.targetPlatform"/>
-							<builder id="com.nvidia.cuda.ide.builder.1635905474" incrementalBuildTarget="${ConfigName}" keepEnvironmentInBuildfile="false" managedBuildOn="false" name="CUDA Toolkit 7.5 Builder" superClass="com.nvidia.cuda.ide.builder"/>
+							<builder buildPath="${workspace_loc:/Project4-CUDA-Rasterizer}/build/" id="com.nvidia.cuda.ide.builder.1635905474" incrementalBuildTarget="${ConfigName}" keepEnvironmentInBuildfile="false" managedBuildOn="false" name="CUDA Toolkit 7.5 Builder" superClass="com.nvidia.cuda.ide.builder"/>
 							<tool id="nvcc.compiler.base.854769442" name="NVCC Compiler" superClass="nvcc.compiler.base">
 								<option id="nvcc.compiler.pic.1870261830" name="Position Independent Code (-fPIC)" superClass="nvcc.compiler.pic"/>
 								<inputType id="nvcc.compiler.input.cu.943232565" superClass="nvcc.compiler.input.cu"/>
@@ -186,9 +186,14 @@
 	</storageModule>
 	<storageModule moduleId="org.eclipse.cdt.core.LanguageSettingsProviders"/>
 	<storageModule moduleId="refreshScope" versionNumber="2">
+		<configuration configurationName="DebugFast"/>
+		<configuration configurationName="DebugNDEBUG"/>
 		<configuration configurationName="Default">
 			<resource resourceType="PROJECT" workspacePath="/Project4-CUDA-Rasterizer"/>
 		</configuration>
+		<configuration configurationName="Debug"/>
+		<configuration configurationName="Release"/>
 	</storageModule>
 	<storageModule moduleId="org.eclipse.cdt.internal.ui.text.commentOwnerProjectMappings"/>
+	<storageModule moduleId="org.eclipse.cdt.make.core.buildtargets"/>
 </cproject>
diff --git a/README.md b/README.md
index cad1abd..342481b 100644
--- a/README.md
+++ b/README.md
@@ -1,20 +1,46 @@
-CUDA Rasterizer
-===============
-
-[CLICK ME FOR INSTRUCTION OF THIS PROJECT](./INSTRUCTION.md)
+CUDA Rasterization Pipeline
+======================
 
 **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)
+* Daniel Krupka
+* Tested on: Debian testing (stretch), Intel(R) Core(TM) i7-4710HQ CPU @ 2.50GHz 8GB, GTX 850M
+
+
+# About
+This is a CUDA rasterization pipeline made for UPenn's CIS565. It features
+* GLTF model loading
+* Perspective correct diffuse PNG textures
+* Screen space ambient occlusion (SSAO)
+The pipeline utilizes deferred shading - the `depthPass` kernel performs depth testing and
+assembles the G-buffer prior to shading and SSAO.
+
+# Screenshots
+Video of the rasterizer in action can be found [here](https://youtu.be/_Y-9eAgICrI).
+![Normals](renders/duck-normals.png "Duck Normals")
+![Lambert](renders/duck-lambert.png "Duck Lambert")
+
 
-### (TODO: Your README)
+# Textures
+I implemented perspective-correct UV texturing for diffuse coloring, using CUDA's texture memory functionality.
+![DuckTex](renders/duck-texture.png "Duck Texture")
+![Truck](renders/truck-nossao.png "truck texture")
 
-*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.
+# SSAO
+I also implemented SSAO, both using shared memory and without.
+![Truck SSAO only](renders/truck-ssao-only.png "Truck SSAO only")
+![Truck SSAO](renders/truck-ssao.png "Truck with ssao")
 
+# Performance
+![notex nosm](renders/plt_loc_notex.png "Duck Texture")
+![tex nosm](renders/plt_loc_tex.png "Duck Texture")
+![tex sm](renders/plt_sm_tex.png "Duck Texture")
+The textureless pipeline was substantially faster for the models with textures (duck, milk truck), particulary when omitting the truck's
+fairly large texture, due to having far fewer memory accesses. However, the small textureless models did experience a small slowdown, possibly due to having
+fewer calculations to mask latency, as barycentric texture coordinate interpolation was skipped. This could be fairly easily extended to include textures for
+bump mapping, specular shading, etc.
 
-### Credits
+For SSAO, shared memory did make a small improvement in all cases, but particularly the cow and duck, which both possessed large overhangs. The `ssaoBlur` kernel
+could have used shared memory as well, but the handling for cross-block contributions was necessary somewhat complicated. Depth testing could also benefit, but
+would require a tiling scheme, also introducing more complication.
 
-* [tinygltfloader](https://github.com/syoyo/tinygltfloader) by [@soyoyo](https://github.com/syoyo)
-* [glTF Sample Models](https://github.com/KhronosGroup/glTF/blob/master/sampleModels/README.md)
diff --git a/cis565_rasterizer.launch b/cis565_rasterizer.launch
index 2c211dc..2d355c1 100644
--- a/cis565_rasterizer.launch
+++ b/cis565_rasterizer.launch
@@ -1,16 +1,24 @@
 <?xml version="1.0" encoding="UTF-8" standalone="no"?>
 <launchConfiguration type="org.eclipse.cdt.launch.applicationLaunchType">
 <stringAttribute key="bad_container_name" value="/Project4-CUDA-Rasterize"/>
+<booleanAttribute key="nvvp.apiTraceProfiling" value="true"/>
+<booleanAttribute key="nvvp.appReplay" value="false"/>
+<booleanAttribute key="nvvp.concurrentProfiling" value="true"/>
+<booleanAttribute key="nvvp.environmentProfiling" value="false"/>
+<intAttribute key="nvvp.multiProcess" value="0"/>
+<booleanAttribute key="nvvp.startProfiling" value="true"/>
+<intAttribute key="nvvp.timeout" value="0"/>
+<booleanAttribute key="nvvp.uvmProfiling" value="true"/>
 <stringAttribute key="org.eclipse.cdt.debug.mi.core.DEBUG_NAME" value="${cuda_tk_bin:/Project4-CUDA-Rasterizer}/cuda-gdb"/>
 <stringAttribute key="org.eclipse.cdt.debug.mi.core.GDB_INIT" value=".cuda-gdbinit"/>
 <stringAttribute key="org.eclipse.cdt.debug.mi.core.protocol" value="mi"/>
 <intAttribute key="org.eclipse.cdt.launch.ATTR_BUILD_BEFORE_LAUNCH_ATTR" value="2"/>
 <stringAttribute key="org.eclipse.cdt.launch.DEBUGGER_ID" value="com.nvidia.cuda.ide.debug.cudagdb"/>
 <stringAttribute key="org.eclipse.cdt.launch.DEBUGGER_START_MODE" value="run"/>
-<stringAttribute key="org.eclipse.cdt.launch.PROGRAM_ARGUMENTS" value="objs/tri.obj"/>
+<stringAttribute key="org.eclipse.cdt.launch.PROGRAM_ARGUMENTS" value="gltfs/box/box.gltf"/>
 <stringAttribute key="org.eclipse.cdt.launch.PROGRAM_NAME" value="build/cis565_rasterizer"/>
 <stringAttribute key="org.eclipse.cdt.launch.PROJECT_ATTR" value="Project4-CUDA-Rasterizer"/>
-<stringAttribute key="org.eclipse.cdt.launch.PROJECT_BUILD_CONFIG_ID_ATTR" value="com.nvidia.cuda.ide.toolchain.on_path.1673801701.1348522543"/>
+<stringAttribute key="org.eclipse.cdt.launch.PROJECT_BUILD_CONFIG_ID_ATTR" value=""/>
 <listAttribute key="org.eclipse.debug.core.MAPPED_RESOURCE_PATHS">
 <listEntry value="/Project4-CUDA-Rasterizer"/>
 </listAttribute>
diff --git a/gltfs/flower/flower.gltf b/gltfs/flower/flower.gltf
index a59e3d2..9023b9d 100644
Binary files a/gltfs/flower/flower.gltf and b/gltfs/flower/flower.gltf differ
diff --git a/renders/2-cylinder-engine-load-correct.png b/renders/2-cylinder-engine-load-correct.png
deleted file mode 100644
index c4c13ce..0000000
Binary files a/renders/2-cylinder-engine-load-correct.png and /dev/null differ
diff --git a/renders/box-shared-bug.png b/renders/box-shared-bug.png
new file mode 100644
index 0000000..26c65a1
Binary files /dev/null and b/renders/box-shared-bug.png differ
diff --git a/renders/duck-diffuse-texture.png b/renders/duck-diffuse-texture.png
deleted file mode 100644
index 26e500e..0000000
Binary files a/renders/duck-diffuse-texture.png and /dev/null differ
diff --git a/renders/duck-lambert.png b/renders/duck-lambert.png
new file mode 100644
index 0000000..f9343e7
Binary files /dev/null and b/renders/duck-lambert.png differ
diff --git a/renders/duck-normals.png b/renders/duck-normals.png
new file mode 100644
index 0000000..0697aca
Binary files /dev/null and b/renders/duck-normals.png differ
diff --git a/renders/duck-texture.png b/renders/duck-texture.png
new file mode 100644
index 0000000..cec7059
Binary files /dev/null and b/renders/duck-texture.png differ
diff --git a/renders/plt_loc_notex.png b/renders/plt_loc_notex.png
new file mode 100644
index 0000000..d022c4c
Binary files /dev/null and b/renders/plt_loc_notex.png differ
diff --git a/renders/plt_loc_tex.png b/renders/plt_loc_tex.png
new file mode 100644
index 0000000..768bcf8
Binary files /dev/null and b/renders/plt_loc_tex.png differ
diff --git a/renders/plt_sm_tex.png b/renders/plt_sm_tex.png
new file mode 100644
index 0000000..1277472
Binary files /dev/null and b/renders/plt_sm_tex.png differ
diff --git a/renders/triangle.png b/renders/triangle.png
deleted file mode 100644
index 95928d2..0000000
Binary files a/renders/triangle.png and /dev/null differ
diff --git a/renders/truck-no-ssao.png b/renders/truck-no-ssao.png
new file mode 100644
index 0000000..51b5323
Binary files /dev/null and b/renders/truck-no-ssao.png differ
diff --git a/renders/truck-ssao-only.png b/renders/truck-ssao-only.png
new file mode 100644
index 0000000..5de9638
Binary files /dev/null and b/renders/truck-ssao-only.png differ
diff --git a/renders/truck-ssao.png b/renders/truck-ssao.png
new file mode 100644
index 0000000..88d5c12
Binary files /dev/null and b/renders/truck-ssao.png differ
diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt
index a57f69f..390f4af 100644
--- 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 -std=c++11 -lineinfo
     )
diff --git a/src/main.cpp b/src/main.cpp
index a36b955..dcfd3ad 100644
--- a/src/main.cpp
+++ b/src/main.cpp
@@ -119,7 +119,7 @@ void runCuda() {
 	glm::mat4 MVP = P * MV;
 
     cudaGLMapBufferObject((void **)&dptr, pbo);
-	rasterize(dptr, MVP, MV, MV_normal);
+	rasterize(dptr, MVP, MV, MV_normal, P);
     cudaGLUnmapBufferObject(pbo);
 
     frame++;
@@ -332,7 +332,7 @@ void keyCallback(GLFWwindow *window, int key, int scancode, int action, int mods
 //----------------------------
 //----- util -----------------
 //----------------------------
-static std::string getFilePathExtension(const std::string &FileName) {
+std::string getFilePathExtension(const std::string &FileName) {
 	if (FileName.find_last_of(".") != std::string::npos)
 		return FileName.substr(FileName.find_last_of(".") + 1);
 	return "";
diff --git a/src/main.hpp b/src/main.hpp
index 4816fa1..edc8ff7 100644
--- 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);
\ No newline at end of file
+void mouseWheelCallback(GLFWwindow* window, double xoffset, double yoffset);
diff --git a/src/rasterize.cu b/src/rasterize.cu
index 4e3504b..52e863d 100644
--- a/src/rasterize.cu
+++ b/src/rasterize.cu
@@ -8,6 +8,7 @@
 
 #include <cmath>
 #include <cstdio>
+#include <cstdint>
 #include <cuda.h>
 #include <cuda_runtime.h>
 #include <thrust/random.h>
@@ -38,15 +39,15 @@ namespace {
 		glm::vec4 pos;
 
 		// TODO: add new attributes to your VertexOut
-		// The attributes listed below might be useful, 
+		// The attributes listed below might be useful,
 		// but always feel free to modify on your own
 
-		 glm::vec3 eyePos;	// eye space position used for shading
-		 glm::vec3 eyeNor;	// eye space normal used for shading, cuz normal will go wrong after perspective transformation
+    glm::vec3 eyePos;	// eye space position used for shading
+    glm::vec3 eyeNor;	// eye space normal used for shading, cuz normal will go wrong after perspective transformation
 		// glm::vec3 col;
-		 glm::vec2 texcoord0;
-		 TextureData* dev_diffuseTex = NULL;
-		// int texWidth, texHeight;
+    glm::vec2 texcoord0;
+    cudaArray* dev_diffuseTex = NULL;
+    cudaTextureObject_t dev_diffuseTexObj = 0;
 		// ...
 	};
 
@@ -56,10 +57,12 @@ namespace {
 	};
 
 	struct Fragment {
-		glm::vec3 color;
+    glm::vec3 eyePos;
+    glm::vec3 eyeNor;
+    float z, ssao;
 
 		// TODO: add new attributes to your Fragment
-		// The attributes listed below might be useful, 
+		// The attributes listed below might be useful,
 		// but always feel free to modify on your own
 
 		// glm::vec3 eyePos;	// eye space position used for shading
@@ -83,7 +86,8 @@ namespace {
 		VertexAttributeTexcoord* dev_texcoord0;
 
 		// Materials, add more attributes when needed
-		TextureData* dev_diffuseTex;
+		cudaArray *dev_diffuseTex;
+    cudaTextureObject_t dev_diffuseTexObj;
 		// TextureData* dev_specularTex;
 		// TextureData* dev_normalTex;
 		// ...
@@ -107,12 +111,14 @@ static Primitive *dev_primitives = NULL;
 static Fragment *dev_fragmentBuffer = NULL;
 static glm::vec3 *dev_framebuffer = NULL;
 
-static int * dev_depth = NULL;	// you might need this buffer when doing depth test
+static unsigned long long *dev_depth = NULL;
+
+static glm::ivec2 *dev_tiles_min = NULL, *dev_tiles_max = NULL;
 
 /**
  * Kernel that writes the image to the OpenGL PBO directly.
  */
-__global__ 
+__global__
 void sendImageToPBO(uchar4 *pbo, int w, int h, glm::vec3 *image) {
     int x = (blockIdx.x * blockDim.x) + threadIdx.x;
     int y = (blockIdx.y * blockDim.y) + threadIdx.y;
@@ -131,53 +137,90 @@ void sendImageToPBO(uchar4 *pbo, int w, int h, glm::vec3 *image) {
     }
 }
 
-/** 
-* Writes fragment colors to the framebuffer
-*/
-__global__
-void render(int w, int h, Fragment *fragmentBuffer, glm::vec3 *framebuffer) {
-    int x = (blockIdx.x * blockDim.x) + threadIdx.x;
-    int y = (blockIdx.y * blockDim.y) + threadIdx.y;
-    int index = x + (y * w);
-
-    if (x < w && y < h) {
-        framebuffer[index] = fragmentBuffer[index].color;
-
-		// TODO: add your fragment shader code here
-
-    }
+static cudaTextureObject_t dev_ssaoTexObj = 0;
+static cudaArray *dev_ssaoTexArray = nullptr;
+static glm::vec3 *dev_ssaoKernel = nullptr;
+static cudaArray *dev_ssaoOutArray = nullptr;
+static cudaSurfaceObject_t dev_ssaoOutSurfObj;
+static cudaTextureObject_t dev_ssaoOutTexObj;
+void ssaoInit(int nKernel, int nRandom, int width, int height) {
+  std::uniform_real_distribution<float> dist(0.0, 1.0);
+  std::default_random_engine rng;
+
+  // generate sampling kernel
+  glm::vec3 *kern = new glm::vec3[nKernel*nKernel];
+  float scale = 1.0f / (nKernel*nKernel);
+  for (int i = 0; i < nKernel*nKernel; i++) {
+    kern[i].x = 2.0f*dist(rng) - 1.0f;
+    kern[i].y = 2.0f*dist(rng) - 1.0f;
+    kern[i].z = dist(rng);
+    kern[i] = glm::normalize(kern[i]) * scale;
+    scale = 0.1f + 0.9f*scale*scale;
+  }
+  cudaMalloc(&dev_ssaoKernel, nKernel*nKernel*sizeof(glm::vec3));
+  cudaMemcpy(dev_ssaoKernel, kern, nKernel*nKernel*sizeof(glm::vec3), cudaMemcpyHostToDevice);
+  delete kern;
+
+  // generate randomization texture
+  float4 *noise = new float4[nRandom*nRandom];
+  for (int i = 0; i < nRandom*nRandom; i++) {
+    noise[i].x = 2.0f*dist(rng) - 1.0f;
+    noise[i].y = 2.0f*dist(rng) - 1.0f;
+    noise[i].z = noise[i].w = 0.0f;
+  }
+  cudaChannelFormatDesc channel = cudaCreateChannelDesc<float4>();
+  cudaMallocArray(&dev_ssaoTexArray, &channel, nRandom, nRandom);
+  cudaMemcpyToArray(dev_ssaoTexArray, 0, 0, noise, nRandom*nRandom*sizeof(float4), cudaMemcpyHostToDevice);
+  checkCUDAError("Set Texture Image data");
+  delete noise;
+
+  // Specify texture
+  cudaResourceDesc resDesc;
+  memset(&resDesc, 0, sizeof(resDesc));
+  resDesc.resType = cudaResourceTypeArray;
+  resDesc.res.array.array = dev_ssaoTexArray;
+
+  // Specify texture object parameters
+  cudaTextureDesc texDesc;
+  memset(&texDesc, 0, sizeof(texDesc));
+  texDesc.addressMode[0]   = cudaAddressModeWrap;
+  texDesc.addressMode[1]   = cudaAddressModeWrap;
+  texDesc.filterMode       = cudaFilterModeLinear;
+  texDesc.readMode         = cudaReadModeElementType;
+  texDesc.normalizedCoords = 1;
+
+  // Create texture object
+  cudaCreateTextureObject(&dev_ssaoTexObj, &resDesc, &texDesc, NULL);
+
+  // Create output array
+  channel = cudaCreateChannelDesc<float>();
+  cudaMallocArray(&dev_ssaoOutArray, &channel, width, height);
+  resDesc.res.array.array = dev_ssaoOutArray;
+  cudaCreateSurfaceObject(&dev_ssaoOutSurfObj, &resDesc);
+  cudaCreateTextureObject(&dev_ssaoOutTexObj, &resDesc, &texDesc, NULL);
 }
 
 /**
  * Called once at the beginning of the program to allocate memory.
  */
 void rasterizeInit(int w, int h) {
-    width = w;
-    height = h;
+  width = w;
+  height = h;
+
 	cudaFree(dev_fragmentBuffer);
 	cudaMalloc(&dev_fragmentBuffer, width * height * sizeof(Fragment));
 	cudaMemset(dev_fragmentBuffer, 0, width * height * sizeof(Fragment));
-    cudaFree(dev_framebuffer);
-    cudaMalloc(&dev_framebuffer,   width * height * sizeof(glm::vec3));
-    cudaMemset(dev_framebuffer, 0, width * height * sizeof(glm::vec3));
-    
-	cudaFree(dev_depth);
-	cudaMalloc(&dev_depth, width * height * sizeof(int));
 
-	checkCUDAError("rasterizeInit");
-}
+  cudaFree(dev_framebuffer);
+  cudaMalloc(&dev_framebuffer,   width * height * sizeof(glm::vec3));
+  cudaMemset(dev_framebuffer, 0, width * height * sizeof(glm::vec3));
 
-__global__
-void initDepth(int w, int h, int * depth)
-{
-	int x = (blockIdx.x * blockDim.x) + threadIdx.x;
-	int y = (blockIdx.y * blockDim.y) + threadIdx.y;
+	cudaFree(dev_depth);
+	cudaMalloc(&dev_depth, width * height * sizeof(unsigned long long));
 
-	if (x < w && y < h)
-	{
-		int index = x + (y * w);
-		depth[index] = INT_MAX;
-	}
+  ssaoInit(8, 4, w, h);
+
+	checkCUDAError("rasterizeInit");
 }
 
 
@@ -185,9 +228,9 @@ void initDepth(int w, int h, int * depth)
 * kern function with support for stride to sometimes replace cudaMemcpy
 * One thread is responsible for copying one component
 */
-__global__ 
+__global__
 void _deviceBufferCopy(int N, BufferByte* dev_dst, const BufferByte* dev_src, int n, int byteStride, int byteOffset, int componentTypeByteSize) {
-	
+
 	// Attribute (vec3 position)
 	// component (3 * float)
 	// byte (4 * byte)
@@ -200,20 +243,20 @@ void _deviceBufferCopy(int N, BufferByte* dev_dst, const BufferByte* dev_src, in
 		int offset = i - count * n;	// which component of the attribute
 
 		for (int j = 0; j < componentTypeByteSize; j++) {
-			
-			dev_dst[count * componentTypeByteSize * n 
-				+ offset * componentTypeByteSize 
+
+			dev_dst[count * componentTypeByteSize * n
+				+ offset * componentTypeByteSize
 				+ j]
 
-				= 
+				=
 
-			dev_src[byteOffset 
-				+ count * (byteStride == 0 ? componentTypeByteSize * n : byteStride) 
-				+ offset * componentTypeByteSize 
+			dev_src[byteOffset
+				+ count * (byteStride == 0 ? componentTypeByteSize * n : byteStride)
+				+ offset * componentTypeByteSize
 				+ j];
 		}
 	}
-	
+
 
 }
 
@@ -233,7 +276,7 @@ void _nodeMatrixTransform(
 }
 
 glm::mat4 getMatrixFromNodeMatrixVector(const tinygltf::Node & n) {
-	
+
 	glm::mat4 curMatrix(1.0);
 
 	const std::vector<double> &m = n.matrix;
@@ -278,7 +321,7 @@ void traverseNode (
 	const tinygltf::Scene & scene,
 	const std::string & nodeString,
 	const glm::mat4 & parentMatrix
-	) 
+	)
 {
 	const tinygltf::Node & n = scene.nodes.at(nodeString);
 	glm::mat4 M = parentMatrix * getMatrixFromNodeMatrixVector(n);
@@ -327,8 +370,8 @@ void rasterizeSetBuffers(const tinygltf::Scene & scene) {
 
 
 
-	// 2. for each mesh: 
-	//		for each primitive: 
+	// 2. for each mesh:
+	//		for each primitive:
 	//			build device buffer of indices, materail, and each attributes
 	//			and store these pointers in a map
 	{
@@ -373,10 +416,10 @@ void rasterizeSetBuffers(const tinygltf::Scene & scene) {
 						return;
 
 					// TODO: add new attributes for your PrimitiveDevBufPointers when you add new attributes
-					VertexIndex* dev_indices;
-					VertexAttributePosition* dev_position;
-					VertexAttributeNormal* dev_normal;
-					VertexAttributeTexcoord* dev_texcoord0;
+					VertexIndex* dev_indices = nullptr;
+					VertexAttributePosition* dev_position = nullptr;
+					VertexAttributeNormal* dev_normal = nullptr;
+					VertexAttributeTexcoord* dev_texcoord0 = nullptr;
 
 					// ----------Indices-------------
 
@@ -516,9 +559,10 @@ void rasterizeSetBuffers(const tinygltf::Scene & scene) {
 
 					// ----------Materials-------------
 
-					// You can only worry about this part once you started to 
+					// You can only worry about this part once you started to
 					// implement textures for your rasterizer
-					TextureData* dev_diffuseTex = NULL;
+					cudaArray* dev_diffuseTex = NULL;
+          cudaTextureObject_t dev_diffuseTexObj = 0;
 					if (!primitive.material.empty()) {
 						const tinygltf::Material &mat = scene.materials.at(primitive.material);
 						printf("material.name = %s\n", mat.name.c_str());
@@ -530,13 +574,38 @@ void rasterizeSetBuffers(const tinygltf::Scene & scene) {
 								if (scene.images.find(tex.source) != scene.images.end()) {
 									const tinygltf::Image &image = scene.images.at(tex.source);
 
-									size_t s = image.image.size() * sizeof(TextureData);
-									cudaMalloc(&dev_diffuseTex, s);
-									cudaMemcpy(dev_diffuseTex, &image.image.at(0), s, cudaMemcpyHostToDevice);
-									
-									// TODO: store the image size to your PrimitiveDevBufPointers
-									// image.width;
-									// image.height;
+                  // convert image to rgba
+                  std::vector<TextureData> rgbaImg;
+                  for (int i = 0; i < image.image.size(); i++) {
+                    rgbaImg.push_back(image.image[i]);
+                    if(i % 3 == 2)
+                      rgbaImg.push_back(0);
+                  }
+
+									size_t s =rgbaImg.size() * sizeof(TextureData);
+                  //printf("img size %d\n", s);
+                  cudaChannelFormatDesc channel = cudaCreateChannelDesc(8,8,8,8,cudaChannelFormatKindUnsigned);
+									cudaMallocArray(&dev_diffuseTex, &channel, image.width, image.height);
+									cudaMemcpyToArray(dev_diffuseTex, 0, 0, &rgbaImg[0], s, cudaMemcpyHostToDevice);
+									checkCUDAError("Set Texture Image data");
+
+                  // Specify texture
+                  cudaResourceDesc resDesc;
+                  memset(&resDesc, 0, sizeof(resDesc));
+                  resDesc.resType = cudaResourceTypeArray;
+                  resDesc.res.array.array = dev_diffuseTex;
+
+                  // Specify texture object parameters
+                  cudaTextureDesc texDesc;
+                  memset(&texDesc, 0, sizeof(texDesc));
+                  texDesc.addressMode[0]   = cudaAddressModeWrap;
+                  texDesc.addressMode[1]   = cudaAddressModeWrap;
+                  texDesc.filterMode       = cudaFilterModeLinear;
+                  texDesc.readMode         = cudaReadModeNormalizedFloat;
+                  texDesc.normalizedCoords = 1;
+
+                  // Create texture object
+                  cudaCreateTextureObject(&dev_diffuseTexObj, &resDesc, &texDesc, NULL);
 
 									checkCUDAError("Set Texture Image data");
 								}
@@ -545,13 +614,13 @@ void rasterizeSetBuffers(const tinygltf::Scene & scene) {
 
 						// TODO: write your code for other materails
 						// You may have to take a look at tinygltfloader
-						// You can also use the above code loading diffuse material as a start point 
+						// You can also use the above code loading diffuse material as a start point
 					}
 
 
 					// ---------Node hierarchy transform--------
 					cudaDeviceSynchronize();
-					
+
 					dim3 numBlocksNodeTransform((numVertices + numThreadsPerBlock.x - 1) / numThreadsPerBlock.x);
 					_nodeMatrixTransform << <numBlocksNodeTransform, numThreadsPerBlock >> > (
 						numVertices,
@@ -577,6 +646,7 @@ void rasterizeSetBuffers(const tinygltf::Scene & scene) {
 						dev_texcoord0,
 
 						dev_diffuseTex,
+            dev_diffuseTexObj,
 
 						dev_vertexOut	//VertexOut
 					});
@@ -590,20 +660,22 @@ void rasterizeSetBuffers(const tinygltf::Scene & scene) {
 		} // for each node
 
 	}
-	
+
 
 	// 3. Malloc for dev_primitives
 	{
 		cudaMalloc(&dev_primitives, totalNumPrimitives * sizeof(Primitive));
+    cudaMalloc(&dev_tiles_min, totalNumPrimitives * sizeof(glm::ivec2));
+    cudaMalloc(&dev_tiles_max, totalNumPrimitives * sizeof(glm::ivec2));
 	}
-	
+
 
 	// Finally, cudaFree raw dev_bufferViews
 	{
 
 		std::map<std::string, BufferByte*>::const_iterator it(bufferViewDevPointers.begin());
 		std::map<std::string, BufferByte*>::const_iterator itEnd(bufferViewDevPointers.end());
-			
+
 			//bufferViewDevPointers
 
 		for (; it != itEnd; it++) {
@@ -616,27 +688,40 @@ void rasterizeSetBuffers(const tinygltf::Scene & scene) {
 
 }
 
-
-
-__global__ 
+__global__
 void _vertexTransformAndAssembly(
-	int numVertices, 
-	PrimitiveDevBufPointers primitive, 
-	glm::mat4 MVP, glm::mat4 MV, glm::mat3 MV_normal, 
+	int numVertices,
+	PrimitiveDevBufPointers primitive,
+	glm::mat4 MVP, glm::mat4 MV, glm::mat3 MV_normal,
 	int width, int height) {
 
 	// vertex id
 	int vid = (blockIdx.x * blockDim.x) + threadIdx.x;
 	if (vid < numVertices) {
 
+    VertexOut &vOut = primitive.dev_verticesOut[vid];
+
 		// TODO: Apply vertex transformation here
 		// Multiply the MVP matrix for each vertex position, this will transform everything into clipping space
 		// Then divide the pos by its w element to transform into NDC space
 		// Finally transform x and y to viewport space
 
-		// TODO: Apply vertex assembly here
-		// Assemble all attribute arraies into the primitive array
-		
+    vOut.pos = MVP*glm::vec4(primitive.dev_position[vid], 1);
+    vOut.pos.x /= vOut.pos.w;
+    vOut.pos.y /= vOut.pos.w;
+    vOut.pos.z /= vOut.pos.w;
+    vOut.pos.x = 0.5f * width * (vOut.pos.x + 1.0f);
+    vOut.pos.y = 0.5f * height * (1.0f - vOut.pos.y);
+
+    vOut.eyePos = glm::vec3(MV*glm::vec4(primitive.dev_position[vid], 1));
+    vOut.eyeNor = glm::normalize(MV_normal*primitive.dev_normal[vid]);
+
+    vOut.dev_diffuseTex = primitive.dev_diffuseTex;
+    vOut.dev_diffuseTexObj = primitive.dev_diffuseTexObj;
+
+    if (primitive.dev_texcoord0)
+      vOut.texcoord0 = primitive.dev_texcoord0[vid];
+
 	}
 }
 
@@ -644,7 +729,7 @@ void _vertexTransformAndAssembly(
 
 static int curPrimitiveBeginId = 0;
 
-__global__ 
+__global__
 void _primitiveAssembly(int numIndices, int curPrimitiveBeginId, Primitive* dev_primitives, PrimitiveDevBufPointers primitive) {
 
 	// index id
@@ -655,78 +740,287 @@ void _primitiveAssembly(int numIndices, int curPrimitiveBeginId, Primitive* dev_
 		// TODO: uncomment the following code for a start
 		// This is primitive assembly for triangles
 
-		//int pid;	// id for cur primitives vector
-		//if (primitive.primitiveMode == TINYGLTF_MODE_TRIANGLES) {
-		//	pid = iid / (int)primitive.primitiveType;
-		//	dev_primitives[pid + curPrimitiveBeginId].v[iid % (int)primitive.primitiveType]
-		//		= primitive.dev_verticesOut[primitive.dev_indices[iid]];
-		//}
-
+		int pid;	// id for cur primitives vector
+		if (primitive.primitiveMode == TINYGLTF_MODE_TRIANGLES) {
+			pid = iid / (int)primitive.primitiveType;
+			dev_primitives[pid + curPrimitiveBeginId].v[iid % (int)primitive.primitiveType]
+				= primitive.dev_verticesOut[primitive.dev_indices[iid]];
+		}
 
 		// TODO: other primitive types (point, line)
 	}
-	
+
 }
 
+// based on http://stereopsis.com/radix.html
+// maps float to uint32 in a way that preserves order. stores in the upper half
+// of a uint64 , and the index of a primitive as the lower 32, so atomicMin can both
+// compare based on depth and store the index
+__device__ static inline uint64_t FloatFlip(float f, int i)
+{
+	uint32_t fi = *((uint32_t*)&f), mask = (fi & 0x80000000) ? 0xFFFFFFFF : 0x80000000;
+	return (((uint64_t)(fi ^ mask)) << 32) | i;
+}
+__device__ static inline float FloatUnflip(uint64_t u)
+{
+  u >>= 32;
+  uint32_t mask = (u & 0x80000000) ? 0xFFFFFFFF : 0x80000000, fi = u ^ mask;
+	return *((float*)&fi);
+}
 
+__global__ void depthPass(int numPrimitives, Primitive *dev_primitives, int w, int h, unsigned long long *depth) {
+  int pIdx = blockIdx.x * blockDim.x + threadIdx.x;
+  if (pIdx >= numPrimitives)
+    return;
+  Primitive &p = dev_primitives[pIdx];
+
+  glm::vec3 tri[3];
+  tri[0] = glm::vec3(p.v[0].pos);
+  tri[1] = glm::vec3(p.v[1].pos);
+  tri[2] = glm::vec3(p.v[2].pos);
+
+  AABB aabb;
+  getAABBForTriangle(tri, aabb);
+  for (int j = aabb.min[1]; j < aabb.max[1] && j < h && j >= 0; j++) {
+  for (int i = aabb.min[0]; i < aabb.max[0] && i < w && i >= 0; i++) {
+    glm::vec2 coord(i,j);
+    glm::vec3 bary = calculateBarycentricCoordinate(tri, coord);
+    if (isBarycentricCoordInBounds(bary))
+      atomicMin(&depth[i+w*j], FloatFlip(getZAtCoordinate(bary, tri),pIdx));
+  }}
+}
+
+
+__device__ void shadeFragment(unsigned int i, unsigned int j, const Primitive &p, glm::vec3 &out) {
+  glm::vec3 tri[3];
+  tri[0] = glm::vec3(p.v[0].pos);
+  tri[1] = glm::vec3(p.v[1].pos);
+  tri[2] = glm::vec3(p.v[2].pos);
+  glm::vec2 coord(i,j);
+  glm::vec3 bary = calculateBarycentricCoordinate(tri, coord);
+
+  // lambert
+  glm::vec3 fragDir = glm::normalize(bary.x*p.v[0].eyePos + bary.y*p.v[1].eyePos + bary.z*p.v[2].eyePos);
+  glm::vec3 fragNrm = glm::normalize(bary.x*p.v[0].eyeNor + bary.y*p.v[1].eyeNor + bary.z*p.v[2].eyeNor);
+  glm::vec3 lambert = glm::clamp(-glm::vec3(glm::dot(fragDir, fragNrm)), 0.0f, 1.0f);
+
+  glm::vec3 texBary = bary / glm::vec3(p.v[0].pos[3], p.v[1].pos[3], p.v[2].pos[3]);
+
+  glm::vec2 st = texBary[0]*p.v[0].texcoord0 + texBary[1]*p.v[1].texcoord0 + texBary[2]*p.v[2].texcoord0;
+  float norm = texBary[0] + texBary[1] + texBary[2];
+  if (p.v[0].dev_diffuseTex) {
+    float4 rgba = tex2D<float4>(p.v[0].dev_diffuseTexObj, st.x / norm, st.y / norm);
+    lambert *= glm::vec3(rgba.x,rgba.y,rgba.z);
+  }
+
+
+  out = lambert;
+}
+
+__global__ void _fragRasterize(int numPrimitives, Primitive *dev_primitives, Fragment *dev_fragments, int w, int h, unsigned long long *depth, glm::vec3 *framebuffer) {
+  int i = blockIdx.x * blockDim.x + threadIdx.x;
+  int j = blockIdx.y * blockDim.y + threadIdx.y;
+  if (i >= w || j >= h)
+    return;
+
+  int pIdx = depth[i+w*j];
+  if (pIdx < 0)
+    return;
+  Primitive &p = dev_primitives[pIdx];
+  Fragment &f = dev_fragments[i+w*j];
+
+  glm::vec3 tri[3];
+  tri[0] = glm::vec3(p.v[0].pos);
+  tri[1] = glm::vec3(p.v[1].pos);
+  tri[2] = glm::vec3(p.v[2].pos);
+  glm::vec2 coord(i,j);
+  glm::vec3 bary = calculateBarycentricCoordinate(tri, coord);
+  f.z = FloatUnflip(depth[i+w*j]);
+
+  // lambert
+  f.eyePos = glm::normalize(bary.x*p.v[0].eyePos
+                          + bary.y*p.v[1].eyePos
+                          + bary.z*p.v[2].eyePos);
+  f.eyeNor = glm::normalize(bary.x*p.v[0].eyeNor
+                          + bary.y*p.v[1].eyeNor
+                          + bary.z*p.v[2].eyeNor);
+  glm::vec3 lambert = glm::clamp(-glm::vec3(glm::dot(f.eyePos, f.eyeNor)), 0.0f, 1.0f);
+
+  glm::vec3 texBary = bary / glm::vec3(p.v[0].pos[3], p.v[1].pos[3], p.v[2].pos[3]);
+  glm::vec2 st0 = texBary[0]*p.v[0].texcoord0
+        + texBary[1]*p.v[1].texcoord0
+        + texBary[2]*p.v[2].texcoord0;
+  st0 /= texBary[0] + texBary[1] + texBary[2];
+
+  if (p.v[0].dev_diffuseTex) {
+    float4 rgba = tex2D<float4>(p.v[0].dev_diffuseTexObj, st0.x, st0.y);
+    lambert *= glm::vec3(rgba.x,rgba.y,rgba.z);
+  }
+  framebuffer[i + w*j] = lambert;
+}
+
+__device__ static inline float smoothstep(float a, float b, float x) {
+  x = (x-a)/(b-a);
+  x = (x < 0.0f) ? 0.0f : ((x > 1.0f) ? 1.0f : x);
+  return x*x*(3.0f - 2.0f*x);
+}
+
+__global__ void ssaoPass(int w, int h, Fragment *dev_fragments, const glm::mat4 P,
+                          int ssaoTexSize, int ssaoTexObj,
+                          int ssaoKernSize, const glm::vec3 *ssaoKern,
+                          float ssaoRadius, unsigned long long *depth, glm::vec3 *framebuffer) {
+  int i = blockIdx.x * blockDim.x + threadIdx.x;
+  int j = blockIdx.y * blockDim.y + threadIdx.y;
+  if (i >= w || j >= h)
+    return;
+  Fragment &f = dev_fragments[i + w*j];
+
+  float4 rVec4 = tex2D<float4>(ssaoTexObj, float(i)/ssaoTexSize, float(j)/ssaoTexSize);
+  glm::vec3 rVec(rVec4.x, rVec4.y, 0.0);
+  glm::vec3 tVec = glm::normalize(rVec - f.eyeNor * glm::dot(rVec, f.eyeNor));
+  glm::vec3 bVec = glm::cross(f.eyeNor, tVec);
+  glm::mat3 TBN(tVec, bVec, f.eyeNor);
+
+  float ssao = 0.0f;
+  for (int k = 0; k < ssaoKernSize*ssaoKernSize; k++) {
+    glm::vec4 samp = P*glm::vec4(ssaoRadius*TBN*ssaoKern[k] + f.eyePos, 1.0f);
+    int si = 0.5f * w * (samp.x + 1.0f);
+    int sj = 0.5f * h * (1.0f - samp.y);
+    float z = dev_fragments[si + w*sj].z;
+    if (z > f.z)
+      ssao += smoothstep(0.0, 1.0, ssaoRadius / fabs(f.z - z));
+  }
+  ssao = 1.0 - ssao/(ssaoKernSize*ssaoKernSize);
+  framebuffer[i+w*j] = glm::vec3(ssao);
+}
+
+__global__ void ssaoPassShared(int w, int h, Fragment *dev_fragments, const glm::mat4 P,
+                          int ssaoTexSize, int ssaoTexObj,
+                          int ssaoKernSize, const glm::vec3 *ssaoKern, float ssaoRadius) {
+  extern __shared__ Fragment sFrag[];
+
+  int i = blockIdx.x * blockDim.x + threadIdx.x;
+  int j = blockIdx.y * blockDim.y + threadIdx.y;
+  if (i >= w || j >= h)
+    return;
+
+  // load the fragment into shared memory
+  int tIdx = threadIdx.x + blockDim.x * threadIdx.y;
+  sFrag[tIdx] = dev_fragments[i + w*j];
+  Fragment &f = sFrag[tIdx];
+
+  float4 rVec4 = tex2D<float4>(ssaoTexObj, float(i)/ssaoTexSize, float(j)/ssaoTexSize);
+  glm::vec3 rVec(rVec4.x, rVec4.y, 0.0);
+  glm::vec3 tVec = glm::normalize(rVec - f.eyeNor * glm::dot(rVec, f.eyeNor));
+  glm::vec3 bVec = glm::cross(f.eyeNor, tVec);
+  glm::mat3 TBN(tVec, bVec, f.eyeNor);
+
+  f.ssao = 0.0;
+  for (int k = 0; k < ssaoKernSize*ssaoKernSize; k++) {
+    glm::vec4 samp = P*glm::vec4(ssaoRadius*TBN*ssaoKern[k] + f.eyePos, 1.0f);
+    int si = glm::floor(0.5f * w * (samp.x + 1.0f));
+    int sj = glm::floor(0.5f * h * (1.0f - samp.y));
+    if (si >= 0 && sj >= 0 && si < w && sj < h) {
+      float z = dev_fragments[si + w*sj].z;
+      if (z > f.z)
+        f.ssao += smoothstep(0.0, 1.0, ssaoRadius / fabs(f.z - z));
+    }
+  }
+  dev_fragments[i+w*j].ssao = 1.0 - f.ssao/(ssaoKernSize*ssaoKernSize);
+}
+
+__global__ void ssaoBlur(int w, int h, Fragment *dev_fragments, int ssaoTexSize, glm::vec3 *framebuffer) {
+  extern __shared__ Fragment sFrag[];
+
+  int i = blockIdx.x * blockDim.x + threadIdx.x;
+  int j = blockIdx.y * blockDim.y + threadIdx.y;
+  if (i >= w || j >= h)
+    return;
+
+  // load the fragment into shared memory
+  int tIdx = threadIdx.x + blockDim.x * threadIdx.y;
+  int idx = i + w*j;
+  sFrag[tIdx] = dev_fragments[idx];
+
+  float ssao = 0.0f;
+  int n = 0;
+  for (int sj = j-ssaoTexSize/2; sj < j+ssaoTexSize/2; sj++) {
+  for (int si = i-ssaoTexSize/2; si < i+ssaoTexSize/2; si++) {
+    if (si >= 0 && sj >= 0 && si < w && sj < h)
+      ssao += dev_fragments[si + w*sj].ssao;
+    else
+      continue;
+    n++;
+  }}
+  ssao /= n;
+  framebuffer[idx] *= ssao;
+}
 
 /**
  * Perform rasterization.
  */
-void rasterize(uchar4 *pbo, const glm::mat4 & MVP, const glm::mat4 & MV, const glm::mat3 MV_normal) {
-    int sideLength2d = 8;
-    dim3 blockSize2d(sideLength2d, sideLength2d);
-    dim3 blockCount2d((width  - 1) / blockSize2d.x + 1,
-		(height - 1) / blockSize2d.y + 1);
+void rasterize(uchar4 *pbo, const glm::mat4 & MVP, const glm::mat4 & MV, const glm::mat3 MV_normal, const glm::mat4 &P) {
+  int sideLength2d = 16;
+  dim3 blockSize2d(sideLength2d, sideLength2d);
+  dim3 blockCount2d((width  - 1) / blockSize2d.x + 1,
+  (height - 1) / blockSize2d.y + 1);
 
 	// Execute your rasterization pipeline here
 	// (See README for rasterization pipeline outline.)
 
 	// Vertex Process & primitive assembly
-	{
-		curPrimitiveBeginId = 0;
-		dim3 numThreadsPerBlock(128);
-
-		auto it = mesh2PrimitivesMap.begin();
-		auto itEnd = mesh2PrimitivesMap.end();
-
-		for (; it != itEnd; ++it) {
-			auto p = (it->second).begin();	// each primitive
-			auto pEnd = (it->second).end();
-			for (; p != pEnd; ++p) {
-				dim3 numBlocksForVertices((p->numVertices + numThreadsPerBlock.x - 1) / numThreadsPerBlock.x);
-				dim3 numBlocksForIndices((p->numIndices + numThreadsPerBlock.x - 1) / numThreadsPerBlock.x);
-
-				_vertexTransformAndAssembly << < numBlocksForVertices, numThreadsPerBlock >> >(p->numVertices, *p, MVP, MV, MV_normal, width, height);
-				checkCUDAError("Vertex Processing");
-				cudaDeviceSynchronize();
-				_primitiveAssembly << < numBlocksForIndices, numThreadsPerBlock >> >
-					(p->numIndices, 
-					curPrimitiveBeginId, 
-					dev_primitives, 
-					*p);
-				checkCUDAError("Primitive Assembly");
-
-				curPrimitiveBeginId += p->numPrimitives;
-			}
-		}
+  curPrimitiveBeginId = 0;
+  dim3 numThreadsPerBlock(128);
+
+  auto it = mesh2PrimitivesMap.begin();
+  auto itEnd = mesh2PrimitivesMap.end();
+
+  for (; it != itEnd; ++it) {
+    auto p = (it->second).begin();	// each primitive
+    auto pEnd = (it->second).end();
+    for (; p != pEnd; ++p) {
+      dim3 numBlocksForVertices((p->numVertices + numThreadsPerBlock.x - 1) / numThreadsPerBlock.x);
+      dim3 numBlocksForIndices((p->numIndices + numThreadsPerBlock.x - 1) / numThreadsPerBlock.x);
+
+      _vertexTransformAndAssembly <<< numBlocksForVertices, numThreadsPerBlock >>>(p->numVertices, *p, MVP, MV, MV_normal, width, height);
+      checkCUDAError("Vertex Processing");
+      cudaDeviceSynchronize();
+      _primitiveAssembly <<< numBlocksForIndices, numThreadsPerBlock >>>
+        (p->numIndices,
+        curPrimitiveBeginId,
+        dev_primitives,
+        *p);
+      checkCUDAError("Primitive Assembly");
+
+      curPrimitiveBeginId += p->numPrimitives;
+    }
+  }
+  checkCUDAError("Vertex Processing and Primitive Assembly");
 
-		checkCUDAError("Vertex Processing and Primitive Assembly");
-	}
-	
 	cudaMemset(dev_fragmentBuffer, 0, width * height * sizeof(Fragment));
-	initDepth << <blockCount2d, blockSize2d >> >(width, height, dev_depth);
-	
-	// TODO: rasterize
+  cudaMemset(dev_depth, 0xFF, width * height * sizeof(unsigned long long));
+  cudaMemset(dev_framebuffer, 0, width*height*sizeof(glm::vec3));
 
+  dim3 blockDim1d(1024);
+  dim3 blockCnt1d((totalNumPrimitives + blockDim1d.x - 1)/blockDim1d.x);
 
+  depthPass<<<blockCnt1d,blockDim1d>>>(totalNumPrimitives, dev_primitives, width, height, dev_depth);
+  checkCUDAError("fragDepthFind");
 
-    // Copy depthbuffer colors into framebuffer
-	render << <blockCount2d, blockSize2d >> >(width, height, dev_fragmentBuffer, dev_framebuffer);
-	checkCUDAError("fragment shader");
-    // Copy framebuffer into OpenGL buffer for OpenGL previewing
-    sendImageToPBO<<<blockCount2d, blockSize2d>>>(pbo, width, height, dev_framebuffer);
-    checkCUDAError("copy render result to pbo");
+  _fragRasterize<<<blockCount2d,blockSize2d>>>(totalNumPrimitives, dev_primitives, dev_fragmentBuffer, width, height, dev_depth, dev_framebuffer);
+  checkCUDAError("fragRasterize");
+
+  //cudaMemset(dev_framebuffer, 0, width*height*sizeof(glm::vec3));
+  int smSize = sideLength2d*sideLength2d*sizeof(Fragment);
+  ssaoPassShared<<<blockCount2d,blockSize2d,smSize>>>(width, height, dev_fragmentBuffer, P, 4, dev_ssaoTexObj, 8, dev_ssaoKernel, 5.0);
+  checkCUDAError("fragRasterize");
+
+  ssaoBlur<<<blockCount2d,blockSize2d,smSize>>>(width, height, dev_fragmentBuffer, 4, dev_framebuffer);
+  checkCUDAError("fragRasterize");
+
+  // Copy framebuffer into OpenGL buffer for OpenGL previewing
+  sendImageToPBO<<<blockCount2d, blockSize2d>>>(pbo, width, height, dev_framebuffer);
+  checkCUDAError("copy render result to pbo");
 }
 
 /**
@@ -734,7 +1028,7 @@ void rasterize(uchar4 *pbo, const glm::mat4 & MVP, const glm::mat4 & MV, const g
  */
 void rasterizeFree() {
 
-    // deconstruct primitives attribute/indices device buffer
+  // deconstruct primitives attribute/indices device buffer
 
 	auto it(mesh2PrimitivesMap.begin());
 	auto itEnd(mesh2PrimitivesMap.end());
@@ -744,28 +1038,29 @@ void rasterizeFree() {
 			cudaFree(p->dev_position);
 			cudaFree(p->dev_normal);
 			cudaFree(p->dev_texcoord0);
-			cudaFree(p->dev_diffuseTex);
+			cudaFreeArray(p->dev_diffuseTex);
 
 			cudaFree(p->dev_verticesOut);
-
-			
-			//TODO: release other attributes and materials
 		}
 	}
 
 	////////////
 
-    cudaFree(dev_primitives);
-    dev_primitives = NULL;
+  cudaFree(dev_primitives);
+  dev_primitives = NULL;
+
+  cudaFree(dev_tiles_min);
+  cudaFree(dev_tiles_max);
+  dev_tiles_min = dev_tiles_max = NULL;
 
 	cudaFree(dev_fragmentBuffer);
 	dev_fragmentBuffer = NULL;
 
-    cudaFree(dev_framebuffer);
-    dev_framebuffer = NULL;
+  cudaFree(dev_framebuffer);
+  dev_framebuffer = NULL;
 
 	cudaFree(dev_depth);
-	dev_depth = NULL;
+	dev_depth = nullptr;
 
-    checkCUDAError("rasterize Free");
+  checkCUDAError("rasterize Free");
 }
diff --git a/src/rasterize.h b/src/rasterize.h
index 560aae9..ab90173 100644
--- a/src/rasterize.h
+++ b/src/rasterize.h
@@ -20,5 +20,5 @@ namespace tinygltf{
 void rasterizeInit(int width, int height);
 void rasterizeSetBuffers(const tinygltf::Scene & scene);
 
-void rasterize(uchar4 *pbo, const glm::mat4 & MVP, const glm::mat4 & MV, const glm::mat3 MV_normal);
+void rasterize(uchar4 *pbo, const glm::mat4 & MVP, const glm::mat4 & MV, const glm::mat3 MV_normal, const glm::mat4 &P);
 void rasterizeFree();
diff --git a/src/rasterizeTools.h b/src/rasterizeTools.h
index 46c701e..a2afbbf 100644
--- a/src/rasterizeTools.h
+++ b/src/rasterizeTools.h
@@ -30,8 +30,7 @@ glm::vec3 multiplyMV(glm::mat4 m, glm::vec4 v) {
  * Finds the axis aligned bounding box for a given triangle.
  */
 __host__ __device__ static
-AABB getAABBForTriangle(const glm::vec3 tri[3]) {
-    AABB aabb;
+void getAABBForTriangle(const glm::vec3 tri[3], AABB &aabb) {
     aabb.min = glm::vec3(
             min(min(tri[0].x, tri[1].x), tri[2].x),
             min(min(tri[0].y, tri[1].y), tri[2].y),
@@ -40,7 +39,6 @@ AABB getAABBForTriangle(const glm::vec3 tri[3]) {
             max(max(tri[0].x, tri[1].x), tri[2].x),
             max(max(tri[0].y, tri[1].y), tri[2].y),
             max(max(tri[0].z, tri[1].z), tri[2].z));
-    return aabb;
 }
 
 // CHECKITOUT
@@ -57,7 +55,7 @@ float calculateSignedArea(const glm::vec3 tri[3]) {
  * Helper function for calculating barycentric coordinates.
  */
 __host__ __device__ static
-float calculateBarycentricCoordinateValue(glm::vec2 a, glm::vec2 b, glm::vec2 c, const glm::vec3 tri[3]) {
+float calculateBarycentricCoordinateValue(const glm::vec2 &a, const glm::vec2 &b, const glm::vec2 &c, const glm::vec3 tri[3]) {
     glm::vec3 baryTri[3];
     baryTri[0] = glm::vec3(a, 0);
     baryTri[1] = glm::vec3(b, 0);
@@ -70,7 +68,7 @@ float calculateBarycentricCoordinateValue(glm::vec2 a, glm::vec2 b, glm::vec2 c,
  * Calculate barycentric coordinates.
  */
 __host__ __device__ static
-glm::vec3 calculateBarycentricCoordinate(const glm::vec3 tri[3], glm::vec2 point) {
+glm::vec3 calculateBarycentricCoordinate(const glm::vec3 tri[3], const glm::vec2& point) {
     float beta  = calculateBarycentricCoordinateValue(glm::vec2(tri[0].x, tri[0].y), point, glm::vec2(tri[2].x, tri[2].y), tri);
     float gamma = calculateBarycentricCoordinateValue(glm::vec2(tri[0].x, tri[0].y), glm::vec2(tri[1].x, tri[1].y), point, tri);
     float alpha = 1.0 - beta - gamma;
@@ -82,7 +80,7 @@ glm::vec3 calculateBarycentricCoordinate(const glm::vec3 tri[3], glm::vec2 point
  * Check if a barycentric coordinate is within the boundaries of a triangle.
  */
 __host__ __device__ static
-bool isBarycentricCoordInBounds(const glm::vec3 barycentricCoord) {
+bool isBarycentricCoordInBounds(const glm::vec3 &barycentricCoord) {
     return barycentricCoord.x >= 0.0 && barycentricCoord.x <= 1.0 &&
            barycentricCoord.y >= 0.0 && barycentricCoord.y <= 1.0 &&
            barycentricCoord.z >= 0.0 && barycentricCoord.z <= 1.0;
@@ -94,8 +92,8 @@ bool isBarycentricCoordInBounds(const glm::vec3 barycentricCoord) {
  * (i.e. depth) on the triangle.
  */
 __host__ __device__ static
-float getZAtCoordinate(const glm::vec3 barycentricCoord, const glm::vec3 tri[3]) {
-    return -(barycentricCoord.x * tri[0].z
+float getZAtCoordinate(const glm::vec3 &barycentricCoord, const glm::vec3 tri[3]) {
+    return (barycentricCoord.x * tri[0].z
            + barycentricCoord.y * tri[1].z
            + barycentricCoord.z * tri[2].z);
 }