diff --git a/README.md b/README.md index f044c82..dd53fd1 100644 --- a/README.md +++ b/README.md @@ -3,11 +3,61 @@ CUDA Denoiser For CUDA Path Tracer **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) +Han Wang + +Tested on: Windows 11, 11th Gen Intel(R) Core(TM) i9-11900H @ 2.50GHz 22GB, GTX 3070 Laptop GPU + +|without denoizer|with denoize| +|:-----:|:-----:| +|| +### Analysis +Firstly, Among the tested output, since the denoize time will be affected by the various variables, I analyzed the denoize in the default condition given to us: +800x800 resolution, 80 filter size. The output is: 4.645Ms + + +Based on my observation, with the denoize feature implemented, the iteration amount needed for reaching the acceptable smooth condition was significantly reduced, leading to more efficient and time-saving results in the process. + +Also, based on my test I got the following graph among the resolution and runtime of denoize: + +|Graph1| +|:-----:| +| + + + +It is easy to see that while the resolution increased, the amount of that denoize needed was increased. It is also easy to understand: there are more pixels needed to estimate, and thus takes more time to denoize. + +Additionally, I get the following graph among the filter size and runtime of denoize: + +|Graph2| +|:-----:| +| + +It is easy to see from the graph, that the runtime needed increased a little bit while the filter size increased. But the increase is rather small. For the small wave in the graph, while the filter size increases, I guess it is because the increased amount is small and even hidden. + + +I recorded the visual output for different filter sizes and get the following graph: +(For comparing, the iterations was set to 100) + + +|0|10|20|30|40| +|:-----:|:-----:|:-----:|:-----:|:-----:| +||||| + +|50|60|70|80|90| +|:-----:|:-----:|:-----:|:-----:|:-----:| +||||| + + +It is easy to see, while the filter size increases, the output becomes gradually clear and easy to see. + + +The next thing is material. I tested the diffuse and reflection, there are no significant runtime differences in denoise between them. In my opinion, the material might affect the basic part of ray tracing, the denoize is likely a post-process approach and won't be affected. + + +In different scenes, the denoizing effect will change based on the condition. For example, in the Cornell-celling scene, since there is more light there, the overall scene is brighter and can easily reach the effect Cornell scene needs higher iterations to get. Thus, it can provide better results with a small filter size. For more complex and less light scenes, we will need more iterations to get acceptable output. Otherwise, we need to increase the filter size. + + -### (TODO: Your README) -*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. diff --git a/img/cornell.2023-10-19_01-52-06z.181samp.png b/img/cornell.2023-10-19_01-52-06z.181samp.png new file mode 100644 index 0000000..22d3f24 Binary files /dev/null and b/img/cornell.2023-10-19_01-52-06z.181samp.png differ diff --git a/img/cornell.2023-10-20_23-54-52z.277samp.png b/img/cornell.2023-10-20_23-54-52z.277samp.png new file mode 100644 index 0000000..5acc3fc Binary files /dev/null and b/img/cornell.2023-10-20_23-54-52z.277samp.png differ diff --git a/img/cornell.2023-10-21_00-28-26z.100samp.png b/img/cornell.2023-10-21_00-28-26z.100samp.png new file mode 100644 index 0000000..1677909 Binary files /dev/null and b/img/cornell.2023-10-21_00-28-26z.100samp.png differ diff --git a/img/cornell.2023-10-21_00-33-12z.100samp.png b/img/cornell.2023-10-21_00-33-12z.100samp.png new file mode 100644 index 0000000..3fa3c1d Binary files /dev/null and b/img/cornell.2023-10-21_00-33-12z.100samp.png differ diff --git a/img/cornell.2023-10-21_00-33-40z.100samp.png b/img/cornell.2023-10-21_00-33-40z.100samp.png new file mode 100644 index 0000000..faa6519 Binary files /dev/null and b/img/cornell.2023-10-21_00-33-40z.100samp.png differ diff --git a/img/cornell.2023-10-21_00-34-03z.100samp.png b/img/cornell.2023-10-21_00-34-03z.100samp.png new file mode 100644 index 0000000..ad951c0 Binary files /dev/null and b/img/cornell.2023-10-21_00-34-03z.100samp.png differ diff --git a/img/cornell.2023-10-21_00-34-22z.100samp.png b/img/cornell.2023-10-21_00-34-22z.100samp.png new file mode 100644 index 0000000..75fee1a Binary files /dev/null and b/img/cornell.2023-10-21_00-34-22z.100samp.png differ diff --git a/img/cornell.2023-10-21_00-34-56z.100samp.png b/img/cornell.2023-10-21_00-34-56z.100samp.png new file mode 100644 index 0000000..024da9d Binary files /dev/null and b/img/cornell.2023-10-21_00-34-56z.100samp.png differ diff --git a/img/cornell.2023-10-21_00-35-21z.100samp.png b/img/cornell.2023-10-21_00-35-21z.100samp.png new file mode 100644 index 0000000..35c5111 Binary files /dev/null and b/img/cornell.2023-10-21_00-35-21z.100samp.png differ diff --git a/img/cornell.2023-10-21_00-35-21z.97samp.png b/img/cornell.2023-10-21_00-35-21z.97samp.png new file mode 100644 index 0000000..27e6b4e Binary files /dev/null and b/img/cornell.2023-10-21_00-35-21z.97samp.png differ diff --git a/img/cornell.2023-10-21_00-35-44z.100samp.png b/img/cornell.2023-10-21_00-35-44z.100samp.png new file mode 100644 index 0000000..383e9a4 Binary files /dev/null and b/img/cornell.2023-10-21_00-35-44z.100samp.png differ diff --git a/img/cornell.2023-10-21_00-36-01z.100samp.png b/img/cornell.2023-10-21_00-36-01z.100samp.png new file mode 100644 index 0000000..74bb51c Binary files /dev/null and b/img/cornell.2023-10-21_00-36-01z.100samp.png differ diff --git a/img/cornell.2023-10-21_00-36-13z.100samp.png b/img/cornell.2023-10-21_00-36-13z.100samp.png new file mode 100644 index 0000000..4e6d65e Binary files /dev/null and b/img/cornell.2023-10-21_00-36-13z.100samp.png differ diff --git a/img/denoize_resolution.png b/img/denoize_resolution.png new file mode 100644 index 0000000..25df870 Binary files /dev/null and b/img/denoize_resolution.png differ diff --git a/img/resolution.png b/img/resolution.png new file mode 100644 index 0000000..b7b73b3 Binary files /dev/null and b/img/resolution.png differ diff --git a/scenes/cornell.txt b/scenes/cornell.txt index 83ff820..ca356d4 100644 --- a/scenes/cornell.txt +++ b/scenes/cornell.txt @@ -52,7 +52,7 @@ EMITTANCE 0 CAMERA RES 800 800 FOVY 45 -ITERATIONS 5000 +ITERATIONS 100 DEPTH 8 FILE cornell EYE 0.0 5 10.5 diff --git a/scenes/cornell1000.txt b/scenes/cornell1000.txt new file mode 100644 index 0000000..9e1a4bf --- /dev/null +++ b/scenes/cornell1000.txt @@ -0,0 +1,117 @@ +// Emissive material (light) +MATERIAL 0 +RGB 1 1 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 5 + +// Diffuse white +MATERIAL 1 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Diffuse red +MATERIAL 2 +RGB .85 .35 .35 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Diffuse green +MATERIAL 3 +RGB .35 .85 .35 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Specular white +MATERIAL 4 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB .98 .98 .98 +REFL 1 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Camera +CAMERA +RES 1000 1000 +FOVY 45 +ITERATIONS 5000 +DEPTH 8 +FILE cornell +EYE 0.0 5 10.5 +LOOKAT 0 5 0 +UP 0 1 0 + + +// Ceiling light +OBJECT 0 +cube +material 0 +TRANS 0 10 0 +ROTAT 0 0 0 +SCALE 3 .3 3 + +// Floor +OBJECT 1 +cube +material 1 +TRANS 0 0 0 +ROTAT 0 0 0 +SCALE 10 .01 10 + +// Ceiling +OBJECT 2 +cube +material 1 +TRANS 0 10 0 +ROTAT 0 0 90 +SCALE .01 10 10 + +// Back wall +OBJECT 3 +cube +material 1 +TRANS 0 5 -5 +ROTAT 0 90 0 +SCALE .01 10 10 + +// Left wall +OBJECT 4 +cube +material 2 +TRANS -5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +// Right wall +OBJECT 5 +cube +material 3 +TRANS 5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +// Sphere +OBJECT 6 +sphere +material 4 +TRANS -1 4 -1 +ROTAT 0 0 0 +SCALE 3 3 3 diff --git a/scenes/cornell1100.txt b/scenes/cornell1100.txt new file mode 100644 index 0000000..2baa8db --- /dev/null +++ b/scenes/cornell1100.txt @@ -0,0 +1,117 @@ +// Emissive material (light) +MATERIAL 0 +RGB 1 1 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 5 + +// Diffuse white +MATERIAL 1 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Diffuse red +MATERIAL 2 +RGB .85 .35 .35 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Diffuse green +MATERIAL 3 +RGB .35 .85 .35 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Specular white +MATERIAL 4 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB .98 .98 .98 +REFL 1 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Camera +CAMERA +RES 1100 1100 +FOVY 45 +ITERATIONS 5000 +DEPTH 8 +FILE cornell +EYE 0.0 5 10.5 +LOOKAT 0 5 0 +UP 0 1 0 + + +// Ceiling light +OBJECT 0 +cube +material 0 +TRANS 0 10 0 +ROTAT 0 0 0 +SCALE 3 .3 3 + +// Floor +OBJECT 1 +cube +material 1 +TRANS 0 0 0 +ROTAT 0 0 0 +SCALE 10 .01 10 + +// Ceiling +OBJECT 2 +cube +material 1 +TRANS 0 10 0 +ROTAT 0 0 90 +SCALE .01 10 10 + +// Back wall +OBJECT 3 +cube +material 1 +TRANS 0 5 -5 +ROTAT 0 90 0 +SCALE .01 10 10 + +// Left wall +OBJECT 4 +cube +material 2 +TRANS -5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +// Right wall +OBJECT 5 +cube +material 3 +TRANS 5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +// Sphere +OBJECT 6 +sphere +material 4 +TRANS -1 4 -1 +ROTAT 0 0 0 +SCALE 3 3 3 diff --git a/scenes/cornell900.txt b/scenes/cornell900.txt new file mode 100644 index 0000000..889fedd --- /dev/null +++ b/scenes/cornell900.txt @@ -0,0 +1,117 @@ +// Emissive material (light) +MATERIAL 0 +RGB 1 1 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 5 + +// Diffuse white +MATERIAL 1 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Diffuse red +MATERIAL 2 +RGB .85 .35 .35 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Diffuse green +MATERIAL 3 +RGB .35 .85 .35 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Specular white +MATERIAL 4 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB .98 .98 .98 +REFL 1 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Camera +CAMERA +RES 900 900 +FOVY 45 +ITERATIONS 5000 +DEPTH 8 +FILE cornell +EYE 0.0 5 10.5 +LOOKAT 0 5 0 +UP 0 1 0 + + +// Ceiling light +OBJECT 0 +cube +material 0 +TRANS 0 10 0 +ROTAT 0 0 0 +SCALE 3 .3 3 + +// Floor +OBJECT 1 +cube +material 1 +TRANS 0 0 0 +ROTAT 0 0 0 +SCALE 10 .01 10 + +// Ceiling +OBJECT 2 +cube +material 1 +TRANS 0 10 0 +ROTAT 0 0 90 +SCALE .01 10 10 + +// Back wall +OBJECT 3 +cube +material 1 +TRANS 0 5 -5 +ROTAT 0 90 0 +SCALE .01 10 10 + +// Left wall +OBJECT 4 +cube +material 2 +TRANS -5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +// Right wall +OBJECT 5 +cube +material 3 +TRANS 5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +// Sphere +OBJECT 6 +sphere +material 4 +TRANS -1 4 -1 +ROTAT 0 0 0 +SCALE 3 3 3 diff --git a/src/main.cpp b/src/main.cpp index 4092ae4..295f883 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -65,6 +65,7 @@ int main(int argc, char** argv) { // Set up camera stuff from loaded path tracer settings iteration = 0; renderState = &scene->state; + Camera &cam = renderState->camera; width = cam.resolution.x; height = cam.resolution.y; @@ -162,13 +163,13 @@ void runCuda() { // execute the kernel int frame = 0; - pathtrace(frame, iteration); + pathtrace(frame, iteration, ui_filterSize, ui_colorWeight, ui_normalWeight, ui_positionWeight, ui_denoise); } if (ui_showGbuffer) { showGBuffer(pbo_dptr); } else { - showImage(pbo_dptr, iteration); + showImage(pbo_dptr, iteration, ui_denoise); } // unmap buffer object diff --git a/src/pathtrace.cu b/src/pathtrace.cu index 23e5f90..cd58cc1 100644 --- a/src/pathtrace.cu +++ b/src/pathtrace.cu @@ -1,4 +1,4 @@ -#include + #include #include #include #include @@ -18,6 +18,10 @@ #define FILENAME (strrchr(__FILE__, '/') ? strrchr(__FILE__, '/') + 1 : __FILE__) #define checkCUDAError(msg) checkCUDAErrorFn(msg, FILENAME, __LINE__) + + + + void checkCUDAErrorFn(const char *msg, const char *file, int line) { #if ERRORCHECK cudaDeviceSynchronize(); @@ -84,15 +88,30 @@ __global__ void gbufferToPBO(uchar4* pbo, glm::ivec2 resolution, GBufferPixel* g static Scene * hst_scene = NULL; static glm::vec3 * dev_image = NULL; +static glm::vec3 * dev_image_new = NULL; +static glm::vec3 * dev_image_old = NULL; +static float* dev_kernel; + static Geom * dev_geoms = NULL; static Material * dev_materials = NULL; static PathSegment * dev_paths = NULL; static ShadeableIntersection * dev_intersections = NULL; static GBufferPixel* dev_gBuffer = NULL; + + +static cudaEvent_t event_start = nullptr; +static cudaEvent_t event_end = nullptr; +float prev_elapsed_time_gpu_milliseconds = 0.f; + // TODO: static variables for device memory, any extra info you need, etc // ... void pathtraceInit(Scene *scene) { + + cudaEventCreate(&event_start); + cudaEventCreate(&event_end); + cudaEventRecord(event_start); + hst_scene = scene; const Camera &cam = hst_scene->state.camera; const int pixelcount = cam.resolution.x * cam.resolution.y; @@ -100,6 +119,15 @@ void pathtraceInit(Scene *scene) { cudaMalloc(&dev_image, pixelcount * sizeof(glm::vec3)); cudaMemset(dev_image, 0, pixelcount * sizeof(glm::vec3)); + cudaMalloc(&dev_image_new, pixelcount * sizeof(glm::vec3)); + cudaMemset(dev_image_new, 0, pixelcount * sizeof(glm::vec3)); + + cudaMalloc(&dev_image_old, pixelcount * sizeof(glm::vec3)); + cudaMemset(dev_image_old, 0, pixelcount * sizeof(glm::vec3)); + + cudaMalloc(&dev_kernel, pixelcount * sizeof(glm::vec3)); + + cudaMalloc(&dev_paths, pixelcount * sizeof(PathSegment)); cudaMalloc(&dev_geoms, scene->geoms.size() * sizeof(Geom)); @@ -114,19 +142,39 @@ void pathtraceInit(Scene *scene) { cudaMalloc(&dev_gBuffer, pixelcount * sizeof(GBufferPixel)); // TODO: initialize any extra device memeory you need + // reference www.eso.org/sci/software/esomidas//doc/user/98NOV/volb/node317.html + std::vector kernelValueList = { 0.0039, 0.0156, 0.0234, 0.0156, 0.0039, + 0.0156, 0.0625, 0.0937, 0.0625, 0.0156, + 0.0234, 0.0937, 0.1406, 0.0937, 0.0234, + 0.0156, 0.0625, 0.0937, 0.0625, 0.0156, + 0.0039, 0.0156, 0.0234, 0.0156, 0.0039}; + + + for (int i = 0; i < kernelValueList.size(); i++) { + cudaMemcpy(dev_kernel + i, &kernelValueList[i], sizeof(float), cudaMemcpyHostToDevice); + } + checkCUDAError("pathtraceInit"); } void pathtraceFree() { cudaFree(dev_image); // no-op if dev_image is null + cudaFree(dev_image_new); + cudaFree(dev_image_old); + cudaFree(dev_kernel); + cudaFree(dev_paths); cudaFree(dev_geoms); cudaFree(dev_materials); cudaFree(dev_intersections); cudaFree(dev_gBuffer); // TODO: clean up any extra device memory you created - + if (event_end != nullptr) { + cudaEventDestroy(event_start); + cudaEventDestroy(event_end); + } + checkCUDAError("pathtraceFree"); } @@ -282,6 +330,8 @@ __global__ void generateGBuffer ( if (idx < num_paths) { gBuffer[idx].t = shadeableIntersections[idx].t; + gBuffer[idx].normal = shadeableIntersections[idx].surfaceNormal; + gBuffer[idx].position = getPointOnRay(pathSegments[idx].ray, shadeableIntersections[idx].t); } } @@ -301,7 +351,61 @@ __global__ void finalGather(int nPaths, glm::vec3 * image, PathSegment * iterati * Wrapper for the __global__ call that sets up the kernel calls and does a ton * of memory management */ -void pathtrace(int frame, int iter) { + +__global__ void denoizeKernelFilter(glm::vec3* img_in, glm::vec3* img_out, GBufferPixel* gbuffer, float c_phi, float n_phi, float p_phi, +Camera cam, float step, float* kernel, int iter) { + + int x = (blockIdx.x * blockDim.x) + threadIdx.x; + int y = (blockIdx.y * blockDim.y) + threadIdx.y; + + int org_idx = y * cam.resolution.x + x; + + glm::vec3 cval = img_in[org_idx]; + glm::vec3 nval = gbuffer[org_idx].normal; + glm::vec3 pval = gbuffer[org_idx].position; + glm::vec3 sum = glm::vec3(0); + step = pow(2, step); + float cum_w = 0.0; + for (int i = -2; i < 3; i++) { + for (int q = -2; q < 3; q++) { + int cur_x = x + q * step; + int cur_y = y + i * step; + if (cur_x > 0 && cur_x < cam.resolution.x - 1 && cur_y>0 && cur_y < cam.resolution.y - 1) { + int cur_idx = cur_y * cam.resolution.x + cur_x; + + glm::vec3 ctemp = img_in[cur_idx]; + glm::vec3 t = (cval - ctemp)/ float(iter); + + float dist2 = dot(t, t); + + + float c_w = max(0.001f,min(exp(-(dist2) / c_phi), 1.0f)); + + glm::vec3 ntemp = gbuffer[cur_idx].normal; + t = nval - ntemp; + dist2 = dot(t, t); + float n_w = min(exp(-(dist2) / n_phi), 1.0f); + + glm::vec3 ptemp = gbuffer[cur_idx].position; + t = pval - ptemp; + dist2 = dot(t, t); + float p_w = min(exp(-(dist2) / p_phi), 1.0f); + + float weight = c_w * n_w * p_w; + int kernel_idx = (i + 2) * 5 + (q + 2); + float new_float = kernel[kernel_idx]; + sum += ctemp * weight * new_float; + cum_w += weight * new_float; + + } + } + } + + img_out[org_idx] = sum / cum_w; +} + + +void pathtrace(int frame, int iter, int filter, float c_phi, float n_phi, float p_phi, bool ui_denoize) { const int traceDepth = hst_scene->state.traceDepth; const Camera &cam = hst_scene->state.camera; const int pixelcount = cam.resolution.x * cam.resolution.y; @@ -356,13 +460,13 @@ void pathtrace(int frame, int iter) { // --- PathSegment Tracing Stage --- // Shoot ray into scene, bounce between objects, push shading chunks - // Empty gbuffer - cudaMemset(dev_gBuffer, 0, pixelcount * sizeof(GBufferPixel)); + // Empty gbuffer + cudaMemset(dev_gBuffer, 0, pixelcount * sizeof(GBufferPixel)); // clean shading chunks cudaMemset(dev_intersections, 0, pixelcount * sizeof(ShadeableIntersection)); - bool iterationComplete = false; + bool iterationComplete = false; while (!iterationComplete) { // tracing @@ -377,34 +481,74 @@ void pathtrace(int frame, int iter) { ); checkCUDAError("trace one bounce"); cudaDeviceSynchronize(); - + if (depth == 0) { generateGBuffer<<>>(num_paths, dev_intersections, dev_paths, dev_gBuffer); } depth++; - shadeSimpleMaterials<<>> ( + shadeSimpleMaterials<<>> ( iter, num_paths, dev_intersections, dev_paths, dev_materials - ); - iterationComplete = depth == traceDepth; + ); + iterationComplete = depth == traceDepth; } // Assemble this iteration and apply it to the image dim3 numBlocksPixels = (pixelcount + blockSize1d - 1) / blockSize1d; finalGather<<>>(num_paths, dev_image, dev_paths); + if (ui_denoize) { + cudaEventRecord(event_start); + + cudaMemcpy(dev_image_old,dev_image , pixelcount * sizeof(glm::vec3), cudaMemcpyDeviceToDevice); + + const Camera& cam = hst_scene->state.camera; + const dim3 blockSize2d(8, 8); + const dim3 blocksPerGrid2d( + (cam.resolution.x + blockSize2d.x - 1) / blockSize2d.x, + (cam.resolution.y + blockSize2d.y - 1) / blockSize2d.y); + + for (int stepsize = 0; stepsize > > (dev_image_old, dev_image_new, dev_gBuffer, c_phi, n_phi, p_phi, cam, stepsize, dev_kernel,iter); + /* + if (stepsize != ceil(log2(filter / 2)) - 1) { + + cudaMemcpy(dev_image_old, dev_image_new, pixelcount * sizeof(glm::vec3), cudaMemcpyDeviceToDevice); + }*/ + std::swap(dev_image_old, dev_image_new); + + } + + cudaMemcpy(hst_scene->state.image.data(), dev_image_new, + pixelcount * sizeof(glm::vec3), cudaMemcpyDeviceToHost); + + cudaEventRecord(event_end); + cudaEventSynchronize(event_end); + + cudaEventElapsedTime(&prev_elapsed_time_gpu_milliseconds, event_start, event_end); + + cout << "Recorded time = " << prev_elapsed_time_gpu_milliseconds << std::endl; + + } + else { + cudaMemcpy(hst_scene->state.image.data(), dev_image, + pixelcount * sizeof(glm::vec3), cudaMemcpyDeviceToHost); + + } + + /////////////////////////////////////////////////////////////////////////// // CHECKITOUT: use dev_image as reference if you want to implement saving denoised images. // Otherwise, screenshots are also acceptable. // Retrieve image from GPU - cudaMemcpy(hst_scene->state.image.data(), dev_image, - pixelcount * sizeof(glm::vec3), cudaMemcpyDeviceToHost); + checkCUDAError("pathtrace"); } @@ -421,7 +565,7 @@ void showGBuffer(uchar4* pbo) { gbufferToPBO<<>>(pbo, cam.resolution, dev_gBuffer); } -void showImage(uchar4* pbo, int iter) { +void showImage(uchar4* pbo, int iter, bool ui_denoise) { const Camera &cam = hst_scene->state.camera; const dim3 blockSize2d(8, 8); const dim3 blocksPerGrid2d( @@ -429,5 +573,11 @@ const Camera &cam = hst_scene->state.camera; (cam.resolution.y + blockSize2d.y - 1) / blockSize2d.y); // Send results to OpenGL buffer for rendering - sendImageToPBO<<>>(pbo, cam.resolution, iter, dev_image); + if (ui_denoise) { + sendImageToPBO << > > (pbo, cam.resolution, iter, dev_image_new); + } + else { + sendImageToPBO << > > (pbo, cam.resolution, iter, dev_image); + } + } diff --git a/src/pathtrace.h b/src/pathtrace.h index 9e12f44..b779af9 100644 --- a/src/pathtrace.h +++ b/src/pathtrace.h @@ -5,6 +5,6 @@ void pathtraceInit(Scene *scene); void pathtraceFree(); -void pathtrace(int frame, int iteration); +void pathtrace(int frame, int iteration, int filter, float c_phi, float n_phi, float p_phi, bool ui_denoize); void showGBuffer(uchar4 *pbo); -void showImage(uchar4 *pbo, int iter); +void showImage(uchar4 *pbo, int iter, bool ui_denoize); diff --git a/src/preview.cpp b/src/preview.cpp index 3ca2718..350c119 100644 --- a/src/preview.cpp +++ b/src/preview.cpp @@ -195,6 +195,7 @@ void drawGui(int windowWidth, int windowHeight) { ImGui_ImplGlfw_NewFrame(); ImGui::NewFrame(); + // Dear imgui define ImVec2 minSize(300.f, 220.f); ImVec2 maxSize((float)windowWidth * 0.5, (float)windowHeight * 0.3); @@ -209,12 +210,12 @@ void drawGui(int windowWidth, int windowHeight) { if (ImGui::IsKeyPressed('H')) { ui_hide = !ui_hide; } - + ImGui::Text("Application average %.3f ms/frame (%.1f FPS)", 1000.0f / ImGui::GetIO().Framerate, ImGui::GetIO().Framerate); ImGui::SliderInt("Iterations", &ui_iterations, 1, startupIterations); - + ImGui::Checkbox("Denoise", &ui_denoise); - ImGui::SliderInt("Filter Size", &ui_filterSize, 0, 100); + ImGui::SliderInt("Filter Size", &ui_filterSize, 0, 500); ImGui::SliderFloat("Color Weight", &ui_colorWeight, 0.0f, 10.0f); ImGui::SliderFloat("Normal Weight", &ui_normalWeight, 0.0f, 10.0f); ImGui::SliderFloat("Position Weight", &ui_positionWeight, 0.0f, 10.0f); diff --git a/src/sceneStructs.h b/src/sceneStructs.h index da7e558..f262832 100644 --- a/src/sceneStructs.h +++ b/src/sceneStructs.h @@ -79,4 +79,8 @@ struct ShadeableIntersection { // What information might be helpful for guiding a denoising filter? struct GBufferPixel { float t; + glm::vec3 position; + glm::vec3 normal; + + };