GBuffer (first ray cache) + basic procedural texture (required for texture mapping)

Author: HummaWhite

README.md

@@ -1,7 +1,7 @@
 CUDA Path Tracer
 ================
 
-**University of Pennsylvania, CIS 565: GPU Programming and Architecture, Project 2**
+**University of Pennsylvania, CIS 565: GPU Programming and Architecture, Project 3**
 
 * Chang Liu
   * [LinkedIn](https://www.linkedin.com/in/chang-liu-0451a6208/)
@@ -11,6 +11,99 @@ CUDA Path Tracer
   - i7-12700 @ 4.90GHz with 16GB RAM
   - RTX 3070 Ti Laptop 8GB
 
+## Representative Outcome
+
 ![](./img/photo_realistic.jpg)
 
-![](./img/photo_realistic_2.jpg)
+<div align="center">"Photoly" realistic!</div>
+
+<div align="center">Rendered at 2400x1800, 3000 spp within 7 minutes </div>
+
+![](./img/aperture_custom.jpg)
+
+<div align="center">Star-shaped bokehs</div>
+
+## Introduction
+
+This is our third project of CIS 565 Fall 2022. In this project, our goal is to implement a GPU-accelerated ray tracer with CUDA. 
+
+
+
+## Features
+
+### Visual
+
+#### Direct Lighting with Multiple Importance Sampling
+
+#### Importance Sampled Skybox (Environment Map)
+
+Tired of "virtual artificial" light sources? Let's introduce some real-world li
+
+#### Physically-Based Materials
+
+#### Normal Map & PBR Texture
+
+#### Physically-Based Camera: Depth of Field & Custom Bokeh Shape
+
+This is really my favorite part of the project.
+
+<div align="center">
+    <img src="./img/aperture_off.jpg" width="49%"/>
+    <img src="./img/aperture.jpg" width="49%"/>
+</div>
+
+
+
+#### Xor-Scrambled Sobol Low Discrepancy Sequence
+
+#### Post Processing
+
+
+
+### Performance
+
+#### Stackless SAH-Constructed Bounding Volume Hierarchy
+
+For ray-scene intersection, I did two levels of optimization.
+
+First, I wrote a SAH-based BVH. SAH, the Surface Area Heuristic is a method to decide how to split a set of bounding volumes 
+
+The second level of optimization
+
+#### Single-Kernel Path Tracing
+
+There is a paper . It had an interesting opinion: instead of 
+
+### Other
+
+#### Streamed Path Tracing Using Stream Compaction
+
+#### First Ray Caching (G-Buffer)
+
+Since I implemented anti-aliasing and physically based camera at the very beginning, when I noticed that there is still a requirement in the basic part, I found it 
+
+## Performance Analysis
+
+### How Much GPU Improves Path Tracing Efficiency
+
+I'm able and confident to answer this question because I have one CPU path tracer from undergrad. 
+
+### Why My Multi-Kernel Streamed Path Tracer Not Faster Than Single-Kernel?
+
+To know how streaming the rays can improve path tracing efficiency, I additionally implemented a single-kernel version of this path tracer.
+
+What got me surprised it wasn't efficient as expected. In some scenes, it was even worse.
+
+Using NSight Compute, I inspected 
+
+In general, it's a tradeoff between thread concurrency and time spent accessing global memory.
+
+### Material Sorting: Why Slower
+
+Or, there is another possibility that BSDF sampling and evaluation is not that time consuming as expected. The bottleneck still lies in traversal of acceleration structure.
+
+Therefore, in my opinion, material sorting is best applied when:
+
+- There are many different materials in the scene
+- Primitives sharing the same material are randomly distributed in many small clusters over the scene space. The clusters' sizes in solid angle are typically less than what a GPU warp can cover
+

src/common.h

@@ -11,8 +11,6 @@
 
 #define CAMERA_PANORAMA false
 
-#define CAMERA_APERTURE_MASK false
-
 struct ToneMapping {
     enum {
         None = 0, Filmic = 1, ACES = 2

src/main.cpp

@@ -297,8 +297,8 @@ void mousePositionCallback(GLFWwindow* window, double xpos, double ypos) {
 
 	if (leftMousePressed) {
 		// compute new camera parameters
-		cam.rotation.x -= (xpos - lastX) / width * 20.f;
-		cam.rotation.y += (ypos - lastY) / height * 20.f;
+		cam.rotation.x -= (xpos - lastX) / width * 40.f;
+		cam.rotation.y += (ypos - lastY) / height * 40.f;
 		cam.rotation.y = glm::clamp(cam.rotation.y, -89.9f, 89.9f);
 		State::camChanged = true;
 	}

src/material.h

@@ -9,6 +9,7 @@
 #define MATERIAL_DIELETRIC_USE_SCHLICK_APPROX false
 
 #define NullTextureId -1
+#define ProcTextureId -2
 #define ProceduralTexId -2
 #define InvalidPdf -1.f
 

src/pathtrace.cu

@@ -73,11 +73,58 @@ static thrust::device_ptr<int> devSegmentMatKeysThr;
 
 static glm::vec3* devGBufferPos = nullptr;
 static glm::vec3* devGBufferNorm = nullptr;
+
+#if ENABLE_GBUFFER
+static Intersection* devGBuffer = nullptr;
+#endif
 
 void InitDataContainer(GuiDataContainer* imGuiData) {
 	guiData = imGuiData;
 }
 
+#if ENABLE_GBUFFER
+#endif
+
+__global__ void renderGBuffer(DevScene* scene, Camera cam, Intersection *GBuffer) {
+	int idx = blockDim.x * blockIdx.x + threadIdx.x;
+	int idy = blockDim.y * blockIdx.y + threadIdx.y;
+	if (idx >= cam.resolution.x || idy >= cam.resolution.y) {
+		return;
+	}
+
+	float aspect = float(cam.resolution.x) / cam.resolution.y;
+	float tanFovY = glm::tan(glm::radians(cam.fov.y));
+	glm::vec2 pixelSize = 1.f / glm::vec2(cam.resolution);
+	glm::vec2 scr = glm::vec2(idx, idy) * pixelSize;
+	glm::vec2 ruv = scr + pixelSize * glm::vec2(.5f);
+	ruv = 1.f - ruv * 2.f;
+
+	glm::vec3 pLens(0.f);
+	glm::vec3 pFocusPlane = glm::vec3(ruv * glm::vec2(aspect, 1.f) * tanFovY, 1.f) * cam.focalDist;
+	glm::vec3 dir = pFocusPlane - pLens;
+
+	Ray ray;
+	ray.direction = glm::normalize(glm::mat3(cam.right, cam.up, cam.view) * dir);
+	ray.origin = cam.position + cam.right * pLens.x + cam.up * pLens.y;
+
+	Intersection intersec;
+	scene->intersect(ray, intersec);
+
+	if (intersec.primId != NullPrimitive) {
+		if (scene->materials[intersec.matId].type == Material::Type::Light) {
+#if SCENE_LIGHT_SINGLE_SIDED
+			if (glm::dot(intersec.norm, ray.direction) < 0.f) {
+				intersec.primId = NullPrimitive;
+			}
+#endif
+		}
+		else {
+			intersec.wo = -ray.direction;
+		}
+	}
+	GBuffer[idy * cam.resolution.x + idx] = intersec;
+}
+
 void pathTraceInit(Scene* scene) {
 	hstScene = scene;
 
@@ -100,17 +147,32 @@ void pathTraceInit(Scene* scene) {
 	cudaMalloc(&devSegmentMatKeys, pixelcount * sizeof(int));
 	devIntersecMatKeysThr = thrust::device_ptr<int>(devIntersecMatKeys);
 	devSegmentMatKeysThr = thrust::device_ptr<int>(devSegmentMatKeys);
-
 	checkCUDAError("pathTraceInit");
+
+#if ENABLE_GBUFFER
+	cudaMalloc(&devGBuffer, pixelcount * sizeof(Intersection));
+	const int BlockSize = 8;
+	dim3 blockSize(BlockSize, BlockSize);
+
+	dim3 blockNum((cam.resolution.x + BlockSize - 1) / BlockSize,
+		(cam.resolution.y + BlockSize - 1) / BlockSize
+	);
+	renderGBuffer<<<blockNum, blockSize>>>(hstScene->devScene, cam, devGBuffer);
+	checkCUDAError("GBuffer");
+	std::cout << "[GBuffer generated]" << std::endl;
+#endif
 }
 
 void pathTraceFree() {
-	cudaFree(devImage);  // no-op if devImage is null
-	cudaFree(devPaths);
-	cudaFree(devTerminatedPaths);
-	cudaFree(devIntersections);
-	cudaFree(devIntersecMatKeys);
-	cudaFree(devSegmentMatKeys);
+	cudaSafeFree(devImage);  // no-op if devImage is null
+	cudaSafeFree(devPaths);
+	cudaSafeFree(devTerminatedPaths);
+	cudaSafeFree(devIntersections);
+	cudaSafeFree(devIntersecMatKeys);
+	cudaSafeFree(devSegmentMatKeys);
+#if ENABLE_GBUFFER
+	cudaSafeFree(devGBuffer);
+#endif
 }
 
 /**
@@ -154,19 +216,10 @@ __device__ Ray sampleCamera(DevScene* scene, const Camera& cam, int x, int y, gl
 	return ray;
 }
 
-/**
-* Generate PathSegments with rays from the camera through the screen into the
-* scene, which is the first bounce of rays.
-*
-* Antialiasing - add rays for sub-pixel sampling
-* motion blur - jitter rays "in time"
-* lens effect - jitter ray origin positions based on a lens
-*/
 __global__ void generateRayFromCamera(
 	DevScene* scene, Camera cam, 
 	int iter, int traceDepth, PathSegment* pathSegments
 ) {
-
 	int x = (blockIdx.x * blockDim.x) + threadIdx.x;
 	int y = (blockIdx.y * blockDim.y) + threadIdx.y;
 
@@ -183,18 +236,13 @@ __global__ void generateRayFromCamera(
 	}
 }
 
-__global__ void previewGBuffer(
-	int iter, 
-	DevScene* scene, Camera cam, 
-	glm::vec3* image, int width, int height,
-	int kind
-) {
+__global__ void previewGBuffer(int iter, DevScene* scene, Camera cam, glm::vec3* image, int kind) {
 	int x = blockDim.x * blockIdx.x + threadIdx.x;
 	int y = blockDim.y * blockIdx.y + threadIdx.y;
-	if (x >= width || y >= height) {
+	if (x >= cam.resolution.x || y >= cam.resolution.y) {
 		return;
 	}
-	int index = y * width + x;
+	int index = y * cam.resolution.x + x;
 	Sampler rng = makeSeededRandomEngine(iter, index, 0, scene->sampleSequence);
 
 	Ray ray = sampleCamera(scene, cam, x, y, sample4D(rng));
@@ -215,9 +263,6 @@ __global__ void previewGBuffer(
 	}
 }
 
-// computeIntersections handles generating ray intersections ONLY.
-// Generating new rays is handled in your shader(s).
-// Feel free to modify the code below.
 __global__ void computeIntersections(
 	int depth,
 	int numPaths,
@@ -226,6 +271,9 @@ __global__ void computeIntersections(
 	Intersection* intersections,
 	int* materialKeys,
 	bool sortMaterial
+#if ENABLE_GBUFFER
+	, Intersection* GBuffer
+#endif
 ) {
 	int pathIdx = blockIdx.x * blockDim.x + threadIdx.x;
 
@@ -235,6 +283,13 @@ __global__ void computeIntersections(
 
 	Intersection intersec;
 	PathSegment segment = pathSegments[pathIdx];
+#if ENABLE_GBUFFER
+	if (depth == 0) {
+		intersections[pathIdx] = GBuffer[pathIdx];
+		return;
+	}
+#endif
+
 #if BVH_DISABLE
 	scene->naiveIntersect(segment.ray, intersec);
 #else
@@ -393,20 +448,15 @@ __global__ void finalGather(int nPaths, glm::vec3* image, PathSegment* iteration
 	}
 }
 
-__global__ void singleKernelPT(
-	int iter, int maxDepth,
-	DevScene* scene,
-	Camera cam,
-	glm::vec3* image, int width, int height
-) {
+__global__ void singleKernelPT(int iter, int maxDepth, DevScene* scene, Camera cam, glm::vec3* image) {
 	int x = blockDim.x * blockIdx.x + threadIdx.x;
 	int y = blockDim.y * blockIdx.y + threadIdx.y;
-	if (x >= width || y >= height) {
+	if (x >= cam.resolution.x || y >= cam.resolution.y) {
 		return;
 	}
 	glm::vec3 accRadiance(0.f);
 
-	int index = y * width + x;
+	int index = y * cam.resolution.x + x;
 	Sampler rng = makeSeededRandomEngine(iter, index, 0, scene->sampleSequence);
 
 	Ray ray = sampleCamera(scene, cam, x, y, sample4D(rng));
@@ -508,13 +558,13 @@ WriteRadiance:
 	image[index] += accRadiance;
 }
 
-__global__ void BVHVisualize(int iter, DevScene* scene, Camera cam, glm::vec3* image, int width, int height) {
+__global__ void BVHVisualize(int iter, DevScene* scene, Camera cam, glm::vec3* image) {
 	int x = blockDim.x * blockIdx.x + threadIdx.x;
 	int y = blockDim.y * blockIdx.y + threadIdx.y;
-	if (x >= width || y >= height) {
+	if (x >= cam.resolution.x || y >= cam.resolution.y) {
 		return;
 	}
-	int index = y * width + x;
+	int index = y * cam.resolution.x + x;
 
 	Sampler rng = makeSeededRandomEngine(iter, index, 0, scene->sampleSequence);
 	Ray ray = sampleCamera(scene, cam, x, y, sample4D(rng));
@@ -577,6 +627,9 @@ void pathTrace(uchar4* pbo, int frame, int iter) {
 			int blockNumIntersec = (numPaths + BlockSizeIntersec - 1) / BlockSizeIntersec;
 			computeIntersections<<<blockNumIntersec, BlockSizeIntersec>>>(
 				depth, numPaths, devPaths, hstScene->devScene, devIntersections, devIntersecMatKeys, Settings::sortMaterial
+#if ENABLE_GBUFFER
+				, devGBuffer
+#endif
 			);
 			checkCUDAError("PT::computeInteractions");
 			cudaDeviceSynchronize();
@@ -627,16 +680,13 @@ void pathTrace(uchar4* pbo, int frame, int iter) {
 		dim3 singlePTBlockSize(BlockSizeSinglePTX, BlockSizeSinglePTY);
 
 		if (Settings::tracer == Tracer::SingleKernel) {
-			singleKernelPT<<<singlePTBlockNum, singlePTBlockSize>>>(
-				iter, Settings::traceDepth, hstScene->devScene, cam, devImage, cam.resolution.x, cam.resolution.y);
+			singleKernelPT<<<singlePTBlockNum, singlePTBlockSize>>>(iter, Settings::traceDepth, hstScene->devScene, cam, devImage);
 		}
 		else if (Settings::tracer == Tracer::BVHVisualize) {
-			BVHVisualize<<<singlePTBlockNum, singlePTBlockSize>>>(
-				iter, hstScene->devScene, cam, devImage, cam.resolution.x, cam.resolution.y);
+			BVHVisualize<<<singlePTBlockNum, singlePTBlockSize>>>(iter, hstScene->devScene, cam, devImage);
 		}
 		else {
-			previewGBuffer<<<singlePTBlockNum, singlePTBlockSize>>>(
-				iter, hstScene->devScene, cam, devImage, cam.resolution.x, cam.resolution.y,
+			previewGBuffer<<<singlePTBlockNum, singlePTBlockSize>>>(iter, hstScene->devScene, cam, devImage,
 				Settings::GBufferPreviewOpt);
 		}
 

src/preview.cpp

@@ -256,7 +256,6 @@ void RenderImGui() {
 
 		glm::vec3 lastPos = cam.position;
 		if (ImGui::DragFloat3("Position", glm::value_ptr(cam.position), .1f)) {
-			cam.rotation += cam.position - lastPos;
 			State::camChanged = true;
 		}