material sorting

Author: HummaWhite

src/common.cpp

@@ -1 +1,5 @@
 #include "common.h"
+
+int Settings::toneMapping = ToneMapping::ACES;
+bool Settings::visualizeBVH = false;
+bool Settings::sortMaterial = false;

src/common.h

@@ -1 +1,24 @@
+#pragma once
 #include <iostream>
+
+#define MESH_DATA_STRUCT_OF_ARRAY false
+#define MESH_DATA_INDEXED false
+
+#define DEV_SCENE_PASS_BY_CONST_MEM false
+
+#define SCENE_LIGHT_SINGLE_SIDED true
+
+#define BVH_DEBUG_VISUALIZATION false
+#define BVH_DISABLE false
+
+struct ToneMapping {
+    enum {
+        None = 0, Filmic = 1, ACES = 2
+    };
+};
+
+struct Settings {
+    static int toneMapping;
+    static bool visualizeBVH;
+    static bool sortMaterial;
+};

src/main.cpp

@@ -159,7 +159,7 @@ void saveImage() {
 		for (int y = 0; y < height; y++) {
 			int index = x + (y * width);
 			glm::vec3 color = renderState->image[index] / samples;
-			switch (ToneMapping::method) {
+			switch (Settings::toneMapping) {
 			case ToneMapping::Filmic:
 				color = Math::filmic(color);
 				break;
@@ -243,7 +243,7 @@ void keyCallback(GLFWwindow* window, int key, int scancode, int action, int mods
 			saveImage();
 			break;
 		case GLFW_KEY_T:
-			ToneMapping::method = (ToneMapping::method + 1) % 3;
+			Settings::toneMapping = (Settings::toneMapping + 1) % 3;
 			break;
 		case GLFW_KEY_SPACE:
 			camChanged = true;

src/pathtrace.cu

@@ -5,6 +5,7 @@
 #include <thrust/random.h>
 #include <thrust/remove.h>
 #include <thrust/device_ptr.h>
+#include <thrust/sort.h>
 
 #include "sceneStructs.h"
 #include "material.h"
@@ -18,10 +19,6 @@
 #include "mathUtil.h"
 #include "sampler.h"
 
-#define BVH_DEBUG_VISUALIZATION false
-
-int ToneMapping::method = ToneMapping::ACES;
-
 //Kernel that writes the image to the OpenGL PBO directly.
 __global__ void sendImageToPBO(uchar4* pbo, glm::ivec2 resolution,
 	int iter, glm::vec3* Image, int toneMapping) {
@@ -63,10 +60,16 @@ static glm::vec3* devImage = nullptr;
 static PathSegment* devPaths = nullptr;
 static PathSegment* devTerminatedPaths = nullptr;
 static Intersection* devIntersections = nullptr;
+static int* devIntersecMatKeys = nullptr;
+static int* devSegmentMatKeys = nullptr;
 // TODO: static variables for device memory, any extra info you need, etc
 // ...
 static thrust::device_ptr<PathSegment> devPathsThr;
 static thrust::device_ptr<PathSegment> devTerminatedPathsThr;
+
+static thrust::device_ptr<Intersection> devIntersectionsThr;
+static thrust::device_ptr<int> devIntersecMatKeysThr;
+static thrust::device_ptr<int> devSegmentMatKeysThr;
 
 void InitDataContainer(GuiDataContainer* imGuiData) {
 	guiData = imGuiData;
@@ -88,6 +91,12 @@ void pathTraceInit(Scene* scene) {
 
 	cudaMalloc(&devIntersections, pixelcount * sizeof(Intersection));
 	cudaMemset(devIntersections, 0, pixelcount * sizeof(Intersection));
+	devIntersectionsThr = thrust::device_ptr<Intersection>(devIntersections);
+
+	cudaMalloc(&devIntersecMatKeys, pixelcount * sizeof(int));
+	cudaMalloc(&devSegmentMatKeys, pixelcount * sizeof(int));
+	devIntersecMatKeysThr = thrust::device_ptr<int>(devIntersecMatKeys);
+	devSegmentMatKeysThr = thrust::device_ptr<int>(devSegmentMatKeys);
 
 	checkCUDAError("pathTraceInit");
 }
@@ -97,6 +106,8 @@ void pathTraceFree() {
 	cudaFree(devPaths);
 	cudaFree(devTerminatedPaths);
 	cudaFree(devIntersections);
+	cudaFree(devIntersecMatKeys);
+	cudaFree(devSegmentMatKeys);
 }
 
 /**
@@ -152,17 +163,23 @@ __global__ void computeIntersections(
 	int numPaths,
 	PathSegment* pathSegments,
 	DevScene* scene,
-	Intersection* intersections
+	Intersection* intersections,
+	int* materialKeys,
+	bool sortMaterial
 ) {
 	int pathIdx = blockIdx.x * blockDim.x + threadIdx.x;
 
-	if (pathIdx < numPaths) {
-#if BVH_DEBUG_VISUALIZATION
-	scene->visualizedIntersect(pathSegments[pathIdx].ray, intersections[pathIdx]);
-#else
+	if (pathIdx >= numPaths) {
+		return;
+	}
+
 	Intersection intersec;
 	PathSegment segment = pathSegments[pathIdx];
+#if BVH_DISABLE
+	scene->naiveIntersect(segment.ray, intersec);
+#else
 	scene->intersect(segment.ray, intersec);
+#endif
 
 	if (intersec.primId != NullPrimitive) {
 		if (scene->devMaterials[intersec.matId].type == Material::Type::Light) {
@@ -176,15 +193,20 @@ __global__ void computeIntersections(
 				// If not first ray, preserve previous sampling information for
 				// MIS calculation
 				intersec.prevPos = segment.ray.origin;
+				intersec.prev = segment.prev;
 			}
 		}
 		else {
 			intersec.wo = -segment.ray.direction;
 		}
+		if (sortMaterial) {
+			materialKeys[pathIdx] = intersec.matId;
+		}
 	}
-	intersections[pathIdx] = intersec;
-#endif
+	else if (sortMaterial) {
+		materialKeys[pathIdx] = -1;
 	}
+	intersections[pathIdx] = intersec;
 }
 
 __global__ void computeTerminatedRays(
@@ -248,17 +270,19 @@ __global__ void pathIntegSampleSurface(
 	glm::vec3 accRadiance(0.f);
 
 	if (material.type == Material::Type::Light) {
+		PrevBSDFSampleInfo prev = intersec.prev;
+
 		glm::vec3 radiance = material.baseColor * material.emittance;
 		if (depth == 0) {
 			accRadiance += radiance;
 		}
-		else if (segment.deltaSample) {
+		else if (prev.deltaSample) {
 			accRadiance += radiance * segment.throughput;
 		}
 		else {
 			float lightPdf = Math::pdfAreaToSolidAngle(Math::luminance(radiance) * scene->sumLightPowerInv,
 				intersec.prevPos, intersec.pos, intersec.norm);
-			float BSDFPdf = segment.BSDFPdf;
+			float BSDFPdf = prev.BSDFPdf;
 			accRadiance += radiance * segment.throughput * Math::powerHeuristic(BSDFPdf, lightPdf);
 		}
 		segment.remainingBounces = 0;
@@ -293,8 +317,7 @@ __global__ void pathIntegSampleSurface(
 			segment.throughput *= sample.bsdf / sample.pdf *
 				(deltaSample ? 1.f : Math::absDot(intersec.norm, sample.dir));
 			segment.ray = makeOffsetedRay(intersec.pos, sample.dir);
-			segment.BSDFPdf = sample.pdf;
-			segment.deltaSample = deltaSample;
+			segment.prev = { sample.pdf, deltaSample };
 			segment.remainingBounces--;
 		}
 	}
@@ -403,14 +426,22 @@ void pathTrace(uchar4* pbo, int frame, int iter) {
 			numPaths,
 			devPaths,
 			hstScene->devScene,
-			devIntersections
+			devIntersections,
+			devIntersecMatKeys,
+			Settings::sortMaterial
 		);
 		checkCUDAError("PT::computeInteractions");
 		cudaDeviceSynchronize();
 
 		// TODO: compare between directly shading the path segments and shading
 		// path segments that have been reshuffled to be contiguous in memory.
 
+		if (Settings::sortMaterial) {
+			cudaMemcpyDevToDev(devSegmentMatKeys, devIntersecMatKeys, numPaths * sizeof(int));
+			thrust::sort_by_key(devIntersecMatKeysThr, devIntersecMatKeysThr + numPaths, devIntersectionsThr);
+			thrust::sort_by_key(devSegmentMatKeysThr, devSegmentMatKeysThr + numPaths, devPathsThr);
+		}
+
 		pathIntegSampleSurface<<<numBlocksPathSegmentTracing, blockSize1D>>>(
 			iter, depth, devPaths, devIntersections, hstScene->devScene, numPaths
 		);
@@ -440,7 +471,7 @@ void pathTrace(uchar4* pbo, int frame, int iter) {
 	///////////////////////////////////////////////////////////////////////////
 
 	// Send results to OpenGL buffer for rendering
-	sendImageToPBO<<<blocksPerGrid2D, blockSize2D>>>(pbo, cam.resolution, iter, devImage, ToneMapping::method);
+	sendImageToPBO<<<blocksPerGrid2D, blockSize2D>>>(pbo, cam.resolution, iter, devImage, Settings::toneMapping);
 
 	// Retrieve image from GPU
 	cudaMemcpy(hstScene->state.image.data(), devImage,

src/pathtrace.h

@@ -3,14 +3,7 @@
 #include <device_launch_parameters.h>
 #include <vector>
 #include "scene.h"
-
-class ToneMapping {
-public:
-    enum {
-        None = 0, Filmic = 1, ACES = 2
-    };
-    static int method;
-};
+#include "common.h"
 
 void InitDataContainer(GuiDataContainer* guiData);
 void pathTraceInit(Scene *scene);

src/preview.cpp

@@ -210,24 +210,18 @@ void RenderImGui()
 	static float f = 0.0f;
 	static int counter = 0;
 
-	ImGui::Begin("Path Tracer Analytics");                  // Create a window called "Hello, world!" and append into it.
-	
-	// LOOK: Un-Comment to check the output window and usage
-	//ImGui::Text("This is some useful text.");               // Display some text (you can use a format strings too)
-	//ImGui::Checkbox("Demo Window", &show_demo_window);      // Edit bools storing our window open/close state
-	//ImGui::Checkbox("Another Window", &show_another_window);
-
-	//ImGui::SliderFloat("float", &f, 0.0f, 1.0f);            // Edit 1 float using a slider from 0.0f to 1.0f
-	//ImGui::ColorEdit3("clear color", (float*)&clear_color); // Edit 3 floats representing a color
-
-	//if (ImGui::Button("Button"))                            // Buttons return true when clicked (most widgets return true when edited/activated)
-	//	counter++;
-	//ImGui::SameLine();
-	//ImGui::Text("counter = %d", counter);
-	ImGui::Text("Traced Depth %d", imguiData->TracedDepth);
-	ImGui::Text("Application average %.3f ms/frame (%.1f FPS)", 1000.0f / ImGui::GetIO().Framerate, ImGui::GetIO().Framerate);
-	ImGui::End();
+	ImGui::Begin("Path Tracer Analytics"); {
+		ImGui::Text("Traced Depth %d", imguiData->TracedDepth);
+		ImGui::Text("Application average %.3f ms/frame (%.1f FPS)", 1000.0f / ImGui::GetIO().Framerate, ImGui::GetIO().Framerate);
 
+		ImGui::End();
+	}
+
+	ImGui::Begin("Options"); {
+		ImGui::Checkbox("Sort material", &Settings::sortMaterial);
+
+		ImGui::End();
+	}
 
 	ImGui::Render();
 	ImGui_ImplOpenGL3_RenderDrawData(ImGui::GetDrawData());

src/preview.h

@@ -1,5 +1,7 @@
 #pragma once
 
+#include "common.h"
+
 extern GLuint pbo;
 
 std::string currentTimeString();

src/scene.h

@@ -15,13 +15,7 @@
 #include "image.h"
 #include "bvh.h"
 #include "sampler.h"
-
-#define MESH_DATA_STRUCT_OF_ARRAY false
-#define MESH_DATA_INDEXED false
-
-#define DEV_SCENE_PASS_BY_CONST_MEM false
-
-#define SCENE_LIGHT_SINGLE_SIDED true
+#include "common.h"
 
 struct MeshData {
     void clear() {
@@ -190,6 +184,24 @@ struct DevScene {
         }
     }
 
+    __device__ bool naiveTestOcclusion(glm::vec3 x, glm::vec3 y) {
+        const float Eps = 1e-4f;
+
+        glm::vec3 dir = y - x;
+        float dist = glm::length(dir);
+        dir /= dist;
+        dist -= Eps;
+
+        Ray ray = makeOffsetedRay(x, dir);
+
+        for (int i = 0; i < (BVHSize + 1) / 2; i++) {
+            if (intersectPrimitive(i, ray, dist)) {
+                return true;
+            }
+        }
+        return false;
+    }
+
     __device__ void intersect(Ray ray, Intersection& intersec) {
         float closestDist = FLT_MAX;
         int closestPrimId = NullPrimitive;
@@ -243,7 +255,6 @@ struct DevScene {
         dist -= Eps;
 
         Ray ray = makeOffsetedRay(x, dir);
-        bool hit = false;
 
         MTBVHNode* nodes = devBVHNodes[getMTBVHId(-ray.direction)];
         int node = 0;
@@ -255,15 +266,17 @@ struct DevScene {
             if (boundHit && boundDist < dist) {
                 int primId = nodes[node].primitiveId;
                 if (primId != NullPrimitive) {
-                    hit |= intersectPrimitive(primId, ray, dist);
+                    if (intersectPrimitive(primId, ray, dist)) {
+                        return true;
+                    }
                 }
                 node++;
             }
             else {
                 node = nodes[node].nextNodeIfMiss;
             }
         }
-        return hit;
+        return false;
     }
 
     __device__ void visualizedIntersect(Ray ray, Intersection& intersec) {
@@ -323,7 +336,12 @@ struct DevScene {
         glm::vec3 v2 = devVertices[primId * 3 + 2];
         glm::vec3 sampled = Math::sampleTriangleUniform(v0, v1, v2, r.z, r.w);
 
-        if (testOcclusion(pos, sampled)) {
+#if BVH_DISABLE
+        bool occ = naiveTestOcclusion(pos, sampled);
+#else
+        bool occ = testOcclusion(pos, sampled);
+#endif
+        if (occ) {
             return InvalidPdf;
         }
         glm::vec3 normal = Math::triangleNormal(v0, v1, v2);