diff --git a/README.md b/README.md index f044c82..ea99263 100644 --- a/README.md +++ b/README.md @@ -1,13 +1,101 @@ -CUDA Denoiser For CUDA Path Tracer +CUDA Denoiser For Path Tracer ================================== -**University of Pennsylvania, CIS 565: GPU Programming and Architecture, Project 4** +* Linda Zhu: [LinkedIn](https://www.linkedin.com/in/lindadaism/) +* Tested on: Windows 11, i7-12800H @ 2.40GHz 16GB, NVIDIA GeForce RTX 3070 Ti (Personal Laptop) -* (TODO) YOUR NAME HERE -* Tested on: (TODO) Windows 22, i7-2222 @ 2.22GHz 22GB, GTX 222 222MB (Moore 2222 Lab) +## Overview +In this project, I implemented a pathtracing denoiser that uses geometry buffers (G-buffers) to guide a smoothing filter. -### (TODO: Your README) +My denoising technique is based on the paper "Edge-Avoiding A-Trous Wavelet Transform for fast Global Illumination Filtering," by Dammertz, Sewtz, Hanika, and Lensch. You can find the paper here: https://jo.dreggn.org/home/2010_atrous.pdf -*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. +Denoisers can help produce a smoother appearance in a pathtraced image with fewer samples-per-pixel/iterations, although the actual improvement often varies from scene-to-scene. +Smoothing an image can be accomplished by blurring pixels - a simple pixel-by-pixel blur filter may sample the color from a pixel's neighbors in the image, weight them by distance, and write the result back into the pixel. +However, just running a simple blur filter on an image often reduces the amount of detail, smoothing sharp edges. This can get worse as the blur filter gets larger, or with more blurring passes. +Fortunately in a 3D scene, we can use per-pixel metrics to help the filter detect and preserve edges (Figure 1). + +| Raw Path-Traced Image | Simple Blur | Denoised Guided by G-Buffers | +|---|---|---| +|![](img/results/raw10Iter.png)|![](img/results/blur10Iter.png)|![](img/results/denoised10Iter.png)| + +*Figure 1. Comparison between different path traced outputs at 10 iterations. Denoiser parameters are configured using slidebars in the GUI.* + +These per-pixel metrics can include scene geometry information (hence G-buffer), such as per-pixel normals and per-pixel positions (Figure 2), as well as surface color or albedo for preserving detail in mapped or procedural textures. For this project I only used per-pixel metrics from the "first bounce" of the ray. + + Per-Pixel Normals | Per-pixel Positions (Normalized) | +|---|--- +|![](img/results/normalMap.png)|![](img/results/positionMap.png) + +*Figure 2. Per-pixel normal map (mapped to 0-1 range) and position map (scaled down).* + +## Visual Quality Analysis +The test scene used for the following results is a standard cornell box with a resolution of 800x800 pixels generated at 50 iterations. For the 3 denoiser weights, we use a `color weight` of 150.0, a `normal weight` of 0.35 and a `position weight` of 0.35. + +### Varying Filter Size +According to Table 2 in the Dammertz et al. paper, 5 * 2#Iter corresponds to the total filtersize. Therefore, varying the filter size basically determines how many times the A-Trous denoiser is applied to the input image. We calculate the number of iterations as the following: +``` + int numIter = filterSize < 5 ? 0 : floor(log2(filterSize / 5.f)); +``` +We loop from zero to `numIter` inclusive so the denoiser would still be applied once even if `numIter` is zero. + +| Filter size = 10 | Filter size = 40 | Filter size = 80 | +|---|---|---| +|![](img/results/denoised50Iter_10filter.png)|![](img/results/denoised50Iter_40filter.png)|![](img/results/denoised50Iter_80filter.png)| + +| Filter size = 120 | Filter size = 160 | Filter size = 200 | +|---|---|---| +|![](img/results/denoised50Iter_120filter.png)|![](img/results/denoised50Iter_160filter.png)|![](img/results/denoised50Iter_200filter.png)| + +As can be seen in the outputs above, increasing the filter size in general improves the smoothness of the scene. The visual quality does seem to scale uniformly with the filter size until filter size hits 160. To my human eyes, there's no obvious improvement between filter size = 160 and 200. I also tried to use the [ImageCompare](https://www.textcompare.org/image/) tool to find the differences (pertaining to transparency, different color, anti-aliasing etc.) in those two images and it showed 0% changes. + +### Different Material Types +Fixing filter size to 80x80 in the denoiser, we add another diffuse sphere in the scene in addition to the specular reflective one, and compare the denoised results to the raw path-traced outputs. + + Raw PT | Denoised | +|---|--- +|![](img/results/raw_diffMat.png)|![](img/results/denoised_diffMat.png) + +Denoising is most effective on diffuse materials where there aren't many fine details to capture. This is reasonable because on diffuse materials, rays can be scattered equally in all directions within the hemisphere of a surface normal, so averaging out its neighboring pixels using some heuristics to contribute to the final color of a pixel is a quick and easy way to "smooth out" the appearance. In fact, with fewer path-tracing iterations (as shown in the comparison below), we can achieve a quite decent visual quality on **diffuse materials** if we tune the weights properly. Considering the point of denoising is to reduce the number of samples-per-pixel/pathtracing iterations needed to generate an acceptably smooth image, our denoiser is sufficient for diffuse materials. + +35 iter, color weight=81, normal weight=0.28, position weight=0.22 | 50 iter, color weight=150, normal weight=0.35, position weight=0.35 | +|---|--- +|![](img/results/denoised_diffuse.png)|![](img/results/denoised_diffMat_diffuse.png) + +However, denoising is less effective on specular materials, especially at lower iteration counts. Even with the integration of the edge-avoiding function, the edges of the cornell box and the boundaries of the yellow sphere in the reflection are still very blurry. This is because for the scene reflected from the surface of the specular sphere, the positions and normals of the corresponding pixels are continuous on the sphere, which defeats the purpose of the edge detection algorithm. To retain the fine details on specular surfaces, we still need to rely on the iteration counts. + +All images denoised: + 50 spp | 100 spp | 500 spp | Image diff between 50 vs. 500 spp +|---|---|---|---| +|![](img/results/cornell_denoised_50samp_spec.png)|![](img/results/cornell_denoised_100samp_spec.png)|![](img/results/cornell_denoised_500samp_spec.png)|![](img/results/textcompare-image-diff.png) + +### Varying Light Source +Since we are using the provided super basic version of path tracer, for this comparison between a small light source and a large light source, we increase the default path-tracing iteration to 500 so that there's at least no fireflies in render of the small light source scene. The denoiser spec remains the same. + + Small Light Source | Large Light Source | +|---|--- +|![](img/results/smallLight500Iter.png)|![](img/results/largeLight500Iter.png) + +The scene with the larger ceiling light produces much better denoised results than the scene with the smaller light. With the same number of iterations, our scene can converge faster with a larger light source because the rays have a higher probability to hit the light source. In the smaller light scene, rays are more likely to miss the light, which leads to a noisier image. Imagine how many iterations it would require for the default cornell box to produce an image with a comparable quality as the scene with the ceiling light. The denoiser is able to produce an image much closer to the expected outcome quickly. + +## Performance Analysis +We measure performance by timing the `pathtrace()` kernel and the `denoise()` kernel, which is only run once at the end of pathtracing. We use `cudaEvents` to record the total execution time of the kernel. For the following measurements, we use the same ceiling light cornell box scene and the same denoiser specifications as noted in the visual analysis part. + +### Total Render Time w/wo Denoiser +Denoising time is indepedent of total path tracing iterations since it only runs once at the end. In the graph below, we can see that total denoising time did not vary much between iteration counts. As the time executing `pathtrace()` increases with the iteration count, we spend lesser and lesser a proportion of the total render time denoising the output image. Given the desirable visual improvement, the performance cost of the denoiser is fully acceptable. + +![](/img/stats/denoiserTime.png) + +### Denoising Time vs. Filter Size +As mentioned in the "Varying Filter Size" part of the visual analysis section, we calculate the number of times we perform denoising passes on the render based on the input filter size. It is expected that the execution time of `denoise()` will increase when filter size increases, since each pixel now needs to consider a higher number of neighboring pixels while approximating the edge-avoiding blur. + +The following tests were performed at a fixed 50 iterations for the path tracer. Note that we only measure the average execution of one pass among the 5 total iterations of denoising. + +![](/img/stats/denoiserFiltersize.png) + +### Denoising at Different Resolutions +Denoising time increases along with the increase of the image resolution. This is expected because a higher resolution means we have to run our filter over more pixels, which will directly impact runtime. + +The following tests were performed at a fixed 50 iterations for the path tracer and a filter size of 80x80 for the denoiser. Here we choose some common image resolution sizes. Note that we only measure the average execution of one pass among the 5 total iterations of denoising. + +![](/img/stats/denoiserRes.png) diff --git a/img/results/bloopers/cornell.2023-10-20_02-10-47z.100samp.png b/img/results/bloopers/cornell.2023-10-20_02-10-47z.100samp.png new file mode 100644 index 0000000..71d1b7f Binary files /dev/null and b/img/results/bloopers/cornell.2023-10-20_02-10-47z.100samp.png differ diff --git a/img/results/bloopers/unnormalizedPosMap.png b/img/results/bloopers/unnormalizedPosMap.png new file mode 100644 index 0000000..68cb555 Binary files /dev/null and b/img/results/bloopers/unnormalizedPosMap.png differ diff --git a/img/results/blur10Iter.png b/img/results/blur10Iter.png new file mode 100644 index 0000000..7fe56b6 Binary files /dev/null and b/img/results/blur10Iter.png differ diff --git a/img/results/cornell_denoised.2023-10-21_00-42-57z.500samp.png b/img/results/cornell_denoised.2023-10-21_00-42-57z.500samp.png new file mode 100644 index 0000000..c7baff7 Binary files /dev/null and b/img/results/cornell_denoised.2023-10-21_00-42-57z.500samp.png differ diff --git a/img/results/cornell_denoised.2023-10-21_02-28-44z.50samp.png b/img/results/cornell_denoised.2023-10-21_02-28-44z.50samp.png new file mode 100644 index 0000000..19f72e3 Binary files /dev/null and b/img/results/cornell_denoised.2023-10-21_02-28-44z.50samp.png differ diff --git a/img/results/cornell_denoised.2023-10-21_02-30-21z.50samp.png b/img/results/cornell_denoised.2023-10-21_02-30-21z.50samp.png new file mode 100644 index 0000000..e250782 Binary files /dev/null and b/img/results/cornell_denoised.2023-10-21_02-30-21z.50samp.png differ diff --git a/img/results/cornell_denoised.2023-10-21_02-30-41z.50samp.png b/img/results/cornell_denoised.2023-10-21_02-30-41z.50samp.png new file mode 100644 index 0000000..e250782 Binary files /dev/null and b/img/results/cornell_denoised.2023-10-21_02-30-41z.50samp.png differ diff --git a/img/results/cornell_denoised.2023-10-21_02-31-57z.50samp.png b/img/results/cornell_denoised.2023-10-21_02-31-57z.50samp.png new file mode 100644 index 0000000..a92e656 Binary files /dev/null and b/img/results/cornell_denoised.2023-10-21_02-31-57z.50samp.png differ diff --git a/img/results/cornell_denoised_100samp.png b/img/results/cornell_denoised_100samp.png new file mode 100644 index 0000000..edc4b66 Binary files /dev/null and b/img/results/cornell_denoised_100samp.png differ diff --git a/img/results/cornell_denoised_100samp_spec.png b/img/results/cornell_denoised_100samp_spec.png new file mode 100644 index 0000000..03f2da1 Binary files /dev/null and b/img/results/cornell_denoised_100samp_spec.png differ diff --git a/img/results/cornell_denoised_200samp.png b/img/results/cornell_denoised_200samp.png new file mode 100644 index 0000000..5baaa63 Binary files /dev/null and b/img/results/cornell_denoised_200samp.png differ diff --git a/img/results/cornell_denoised_200samp_spec.png b/img/results/cornell_denoised_200samp_spec.png new file mode 100644 index 0000000..2d0665d Binary files /dev/null and b/img/results/cornell_denoised_200samp_spec.png differ diff --git a/img/results/cornell_denoised_500samp_spec.png b/img/results/cornell_denoised_500samp_spec.png new file mode 100644 index 0000000..3a6af99 Binary files /dev/null and b/img/results/cornell_denoised_500samp_spec.png differ diff --git a/img/results/cornell_denoised_50samp_spec.png b/img/results/cornell_denoised_50samp_spec.png new file mode 100644 index 0000000..1a6454f Binary files /dev/null and b/img/results/cornell_denoised_50samp_spec.png differ diff --git a/img/results/denoised10Iter.png b/img/results/denoised10Iter.png new file mode 100644 index 0000000..d287cb0 Binary files /dev/null and b/img/results/denoised10Iter.png differ diff --git a/img/results/denoised50Iter_10filter.png b/img/results/denoised50Iter_10filter.png new file mode 100644 index 0000000..3cb6b02 Binary files /dev/null and b/img/results/denoised50Iter_10filter.png differ diff --git a/img/results/denoised50Iter_120filter.png b/img/results/denoised50Iter_120filter.png new file mode 100644 index 0000000..07eedc0 Binary files /dev/null and b/img/results/denoised50Iter_120filter.png differ diff --git a/img/results/denoised50Iter_160filter.png b/img/results/denoised50Iter_160filter.png new file mode 100644 index 0000000..320e43e Binary files /dev/null and b/img/results/denoised50Iter_160filter.png differ diff --git a/img/results/denoised50Iter_200filter.png b/img/results/denoised50Iter_200filter.png new file mode 100644 index 0000000..320e43e Binary files /dev/null and b/img/results/denoised50Iter_200filter.png differ diff --git a/img/results/denoised50Iter_40filter.png b/img/results/denoised50Iter_40filter.png new file mode 100644 index 0000000..5db4c92 Binary files /dev/null and b/img/results/denoised50Iter_40filter.png differ diff --git a/img/results/denoised50Iter_80filter.png b/img/results/denoised50Iter_80filter.png new file mode 100644 index 0000000..07eedc0 Binary files /dev/null and b/img/results/denoised50Iter_80filter.png differ diff --git a/img/results/denoised_diffMat.png b/img/results/denoised_diffMat.png new file mode 100644 index 0000000..ddecfc5 Binary files /dev/null and b/img/results/denoised_diffMat.png differ diff --git a/img/results/denoised_diffMat_diffuse.png b/img/results/denoised_diffMat_diffuse.png new file mode 100644 index 0000000..a7560f6 Binary files /dev/null and b/img/results/denoised_diffMat_diffuse.png differ diff --git a/img/results/denoised_diffuse.png b/img/results/denoised_diffuse.png new file mode 100644 index 0000000..e0dd83b Binary files /dev/null and b/img/results/denoised_diffuse.png differ diff --git a/img/results/halfNormalizedPosMap.png b/img/results/halfNormalizedPosMap.png new file mode 100644 index 0000000..fc4bf30 Binary files /dev/null and b/img/results/halfNormalizedPosMap.png differ diff --git a/img/results/largeLight500Iter.png b/img/results/largeLight500Iter.png new file mode 100644 index 0000000..c823d96 Binary files /dev/null and b/img/results/largeLight500Iter.png differ diff --git a/img/results/normalMap.png b/img/results/normalMap.png new file mode 100644 index 0000000..e56171d Binary files /dev/null and b/img/results/normalMap.png differ diff --git a/img/results/positionMap.png b/img/results/positionMap.png new file mode 100644 index 0000000..88f2ce5 Binary files /dev/null and b/img/results/positionMap.png differ diff --git a/img/results/raw10Iter.png b/img/results/raw10Iter.png new file mode 100644 index 0000000..c9497f8 Binary files /dev/null and b/img/results/raw10Iter.png differ diff --git a/img/results/raw50Iter.png b/img/results/raw50Iter.png new file mode 100644 index 0000000..7e7f18c Binary files /dev/null and b/img/results/raw50Iter.png differ diff --git a/img/results/raw_diffMat.png b/img/results/raw_diffMat.png new file mode 100644 index 0000000..0469953 Binary files /dev/null and b/img/results/raw_diffMat.png differ diff --git a/img/results/smallLight500Iter.png b/img/results/smallLight500Iter.png new file mode 100644 index 0000000..25c8b60 Binary files /dev/null and b/img/results/smallLight500Iter.png differ diff --git a/img/results/textcompare-image-diff.png b/img/results/textcompare-image-diff.png new file mode 100644 index 0000000..f3b1042 Binary files /dev/null and b/img/results/textcompare-image-diff.png differ diff --git a/img/stats/denoiserFiltersize.png b/img/stats/denoiserFiltersize.png new file mode 100644 index 0000000..154471b Binary files /dev/null and b/img/stats/denoiserFiltersize.png differ diff --git a/img/stats/denoiserRes.png b/img/stats/denoiserRes.png new file mode 100644 index 0000000..80c156b Binary files /dev/null and b/img/stats/denoiserRes.png differ diff --git a/img/stats/denoiserTime.png b/img/stats/denoiserTime.png new file mode 100644 index 0000000..ea4e8c7 Binary files /dev/null and b/img/stats/denoiserTime.png differ diff --git a/scenes/cornell.txt b/scenes/cornell.txt index 83ff820..82cd9b1 100644 --- a/scenes/cornell.txt +++ b/scenes/cornell.txt @@ -52,7 +52,7 @@ EMITTANCE 0 CAMERA RES 800 800 FOVY 45 -ITERATIONS 5000 +ITERATIONS 500 DEPTH 8 FILE cornell EYE 0.0 5 10.5 diff --git a/scenes/cornell_ceiling_light.txt b/scenes/cornell_ceiling_light.txt index 15af5f1..cd31306 100644 --- a/scenes/cornell_ceiling_light.txt +++ b/scenes/cornell_ceiling_light.txt @@ -48,11 +48,51 @@ REFR 0 REFRIOR 0 EMITTANCE 0 +// Refractive white +MATERIAL 5 +RGB 0 0 0 +SPECEX 0 +SPECRGB .98 .98 .98 +REFL 0 +REFR 1 +REFRIOR 1.55 +EMITTANCE 0 + +// Glass white +MATERIAL 6 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB .98 .98 .98 +REFL 1 +REFR 1 +REFRIOR 1.55 +EMITTANCE 0 + +// Plastic white but turns out black?? +MATERIAL 7 +RGB 0.98 0.98 0.98 +SPECEX 0 +SPECRGB 0.98 0.98 0.98 +REFL 1 +REFR 0 +REFRIOR 1.55 +EMITTANCE 0 + +// Diffuse blue +MATERIAL 8 +RGB 1 0.98 0.01 +SPECEX 0 +SPECRGB 0.98 0.98 0.98 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + // Camera CAMERA -RES 800 800 +RES 2560 2560 FOVY 45 -ITERATIONS 10 +ITERATIONS 500 DEPTH 8 FILE cornell EYE 0.0 5 10.5 @@ -115,3 +155,12 @@ material 4 TRANS -1 4 -1 ROTAT 0 0 0 SCALE 3 3 3 + +// Sphere +//OBJECT 7 +//sphere +//material 8 +//TRANS 2.5 2 1 +//ROTAT 0 0 0 +//SCALE 3 3 3 + diff --git a/src/interactions.h b/src/interactions.h index 144a9f5..bf57178 100644 --- a/src/interactions.h +++ b/src/interactions.h @@ -40,23 +40,221 @@ glm::vec3 calculateRandomDirectionInHemisphere( + sin(around) * over * perpendicularDirection2; } +__host__ __device__ +glm::vec3 sampleDiffuse(const Material& m, glm::vec3 nor, + thrust::default_random_engine& rng, glm::vec3& wi) { // write out to wi + + wi = calculateRandomDirectionInHemisphere(nor, rng); + return m.color; +} + +__host__ __device__ +glm::vec3 sampleDiffuseTrans(const Material& m, glm::vec3 nor, + thrust::default_random_engine& rng, glm::vec3& wi, + float& absDot, float& pdf) { + // sample wi in the opposite hemisphere of wo, everything else same as diffuse_refl + wi = calculateRandomDirectionInHemisphere(nor, rng); + wi.z = -wi.z; + absDot = glm::abs(glm::dot(nor, wi)); + pdf = absDot * INV_PI; + return m.color * INV_PI; +} + +__host__ __device__ +glm::vec3 sampleSpecularRefl(const Material& m, glm::vec3 nor, + glm::vec3 wo, glm::vec3& wi) { + + wi = glm::reflect(wo, nor); + return m.specular.color; +} + +__host__ __device__ +glm::vec3 sampleSpecularTrans(const Material& m, glm::vec3 nor, + glm::vec3 wo, glm::vec3& wi) { + + // Default to hard-coded glass IOR if not provided + float ior = m.indexOfRefraction; + ior = ior ? 1.55f : ior; + + // figure out which surface to enter/exit + // eta of air always assumed to be 1 + bool entering = glm::clamp(glm::dot(wo, nor), -1.f, 1.f) < 0.0f; + float etaI = entering ? 1.0f : ior; // incident index + float etaT = entering ? ior : 1.0f; // transmitted index + float eta = etaI / etaT; + + // flip normal + nor = entering ? nor : -nor; + wi = glm::refract(wo, nor, eta); + + // handle total internal refl --> basically specular BRDF + if (glm::length(wi) < EPSILON) { + wi = glm::reflect(wo, nor); + return glm::vec3(0.0f); + } + return m.specular.color; +} + +__host__ __device__ +float fresnelDielectricEval(const Material& m, float cosThetaI) { + float etaI = 1.f; + float etaT = m.indexOfRefraction; + etaT = etaT < EPSILON ? 1.55f : etaT; + cosThetaI = glm::clamp(cosThetaI, -1.f, 1.f); + + // see pbrt FrDielectric() + // Potentially swap indices of refraction + if (cosThetaI > 0.f) { + float temp = etaI; + etaI = etaT; + etaT = temp; + } + cosThetaI = glm::abs(cosThetaI); + + // Computer cosThetaT using Snell's law + float sinThetaI = glm::sqrt(glm::max(0.f, + 1.f - cosThetaI * cosThetaI)); + float sinThetaT = etaI / etaT * sinThetaI; + + // Handle total internal reflection + if (sinThetaT >= 0.999f) { + return 1.f; + } + + // Compute Fresnel reflectance using light polarization eqns, see PBRT 8.2.1 + float cosThetaT = glm::sqrt(glm::max(0.f, + 1.f - sinThetaT * sinThetaT)); + float Rparl = ((etaT * cosThetaI) - (etaI * cosThetaT)) / + ((etaT * cosThetaI) + (etaI * cosThetaT)); + float Rperp = ((etaI * cosThetaI) - (etaT * cosThetaT)) / + ((etaI * cosThetaI) + (etaT * cosThetaT)); + + return (Rparl * Rparl + Rperp * Rperp) * 0.5f; // coefficient +} + +__host__ __device__ +glm::vec3 sampleGlass(const Material& m, glm::vec3 nor, + thrust::default_random_engine& rng, glm::vec3 wo, glm::vec3& wi, + float& absDot, float& pdf) { + + thrust::uniform_real_distribution u01(0, 1); + bool random = u01(rng); + + float fresnel = fresnelDielectricEval(m, glm::dot(wo, nor)); + glm::vec3 bsdf(0.f); + if (random < fresnel) { + // Have to double contribution b/c we only sample + // reflection BxDF half the time + bsdf = sampleSpecularRefl(m, nor, wo, wi); + return bsdf; + } + else { + bsdf = sampleSpecularTrans(m, nor, wo, wi); + return bsdf; // 1-fr b/c all conditions sum up to 1 + } +} + +__host__ __device__ +glm::vec3 samplePlastic(const Material& m, glm::vec3 nor, + thrust::default_random_engine& rng, glm::vec3 wo, glm::vec3& wi, + float& absDot, float& pdf) { + + thrust::uniform_real_distribution u01(0, 1); + bool random = u01(rng); + glm::vec3 bsdf(0.f); + + float fresnel = fresnelDielectricEval(m, glm::dot(wo, nor));; + //if (random < fresnel) { + // // Have to double contribution b/c we only sample + // // reflection BxDF half the time + // bsdf = sampleSpecularRefl(m, nor, wo, wi); + // return bsdf; + //} + //else { + // bsdf = sampleDiffuse(m, nor, rng, wi); + // return bsdf; // 1-fr b/c all conditions sum up to 1 + //} + + if (random < 0.5f) { + // diffuse + bsdf = sampleDiffuse(m, nor, rng, wi);// glm::normalize(calculateRandomDirectionInHemisphere(nor, rng)); + absDot = glm::abs(glm::dot(nor, wi)); + pdf = absDot * INV_PI; + bsdf = m.color * INV_PI; + bsdf *= (1.f - fresnel); + } + else { + // spec refl + bsdf = sampleSpecularRefl(m, nor, wo, wi);// glm::reflect(wo, nor); + absDot = glm::abs(glm::dot(nor, wi)); + pdf = 1.0f; + if (absDot > 0.0f) { + bsdf = m.specular.color / absDot; + } + bsdf *= fresnel; + } + return bsdf * 2.0f; +} + /** * Simple ray scattering with diffuse and perfect specular support. */ __host__ __device__ void scatterRay( - PathSegment & pathSegment, - glm::vec3 intersect, - glm::vec3 normal, - const Material &m, - thrust::default_random_engine &rng) { + PathSegment& pathSegment, + glm::vec3 intersect, + glm::vec3 normal, + const Material& m, + thrust::default_random_engine& rng) { glm::vec3 newDirection; if (m.hasReflective) { newDirection = glm::reflect(pathSegment.ray.direction, normal); - } else { + } + else { newDirection = calculateRandomDirectionInHemisphere(normal, rng); } pathSegment.ray.direction = newDirection; pathSegment.ray.origin = intersect + (newDirection * 0.0001f); } + +__host__ __device__ +void scatterRayMoreMats( + PathSegment & pathSegment, + glm::vec3 intersect, + glm::vec3 normal, + const Material &m, + thrust::default_random_engine &rng) { + if (pathSegment.remainingBounces == 0) return; + + glm::vec3 wo = pathSegment.ray.direction; + glm::vec3 wi(0.0f); + glm::vec3 bsdf(0.0f); + float absDot = 1.f, pdf = 1.f; + bool isDiffuse = false; + bool manualCalc = false; + + // Based on material type + if (m.hasReflective && m.hasRefractive) { + bsdf = sampleGlass(m, normal, rng, wo, wi, absDot, pdf); + } + // Had to manually calculate absDot and pdf to get correct results, otherwise not working properly + else if (m.hasReflective && glm::length(m.color) > EPSILON) { + bsdf = samplePlastic(m, normal, rng, wo, wi, absDot, pdf); + manualCalc = true; + } + else if (m.hasReflective) { + bsdf = sampleSpecularRefl(m, normal, wo, wi); + } + else if (m.hasRefractive) { + bsdf = sampleSpecularTrans(m, normal, wo, wi); + } + else { // default to lambert diffuse + bsdf = sampleDiffuse(m, normal, rng, wi); + isDiffuse = true; + } + + pathSegment.ray.direction = glm::normalize(wi); + pathSegment.ray.origin = intersect + (glm::normalize(wi) * 0.0001f); + pathSegment.throughput = manualCalc ? (bsdf * absDot / pdf) : bsdf; +} diff --git a/src/main.cpp b/src/main.cpp index 4092ae4..e0b3a9b 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -7,6 +7,7 @@ #include "../imgui/imgui_impl_opengl3.h" static std::string startTimeString; +static std::string resultsFolderPath = "../img/results/"; // For camera controls static bool leftMousePressed = false; @@ -23,12 +24,16 @@ int ui_iterations = 0; int startupIterations = 0; int lastLoopIterations = 0; bool ui_showGbuffer = false; +int ui_currGBuffer = 0; bool ui_denoise = false; +bool ui_atrous = true; int ui_filterSize = 80; -float ui_colorWeight = 0.45f; +float ui_colorWeight = 150.f; float ui_normalWeight = 0.35f; -float ui_positionWeight = 0.2f; +float ui_positionWeight = 0.35f; bool ui_saveAndExit = false; +int denoiserCallCount = 0; +int ptCallCount = 0; static bool camchanged = true; static float dtheta = 0, dphi = 0; @@ -69,8 +74,8 @@ int main(int argc, char** argv) { width = cam.resolution.x; height = cam.resolution.y; - ui_iterations = renderState->iterations; - startupIterations = ui_iterations; + ui_iterations = 50; // renderState->iterations; + startupIterations = renderState->iterations; glm::vec3 view = cam.view; glm::vec3 up = cam.up; @@ -111,8 +116,12 @@ void saveImage() { } std::string filename = renderState->imageName; + if (ui_denoise) { + filename += "_denoised"; + } + std::ostringstream ss; - ss << filename << "." << startTimeString << "." << samples << "samp"; + ss << resultsFolderPath << filename << "." << startTimeString << "." << samples << "samp"; filename = ss.str(); // CHECKITOUT @@ -128,6 +137,10 @@ void runCuda() { if (camchanged) { iteration = 0; + // for performance analysis + denoiserCallCount = 0; + ptCallCount = 0; + Camera &cam = renderState->camera; cameraPosition.x = zoom * sin(phi) * sin(theta); cameraPosition.y = zoom * cos(theta); @@ -157,18 +170,37 @@ void runCuda() { uchar4 *pbo_dptr = NULL; cudaGLMapBufferObject((void**)&pbo_dptr, pbo); + PerformanceTimer timer; + timer.startCpuTimer(); + if (iteration < ui_iterations) { iteration++; - // execute the kernel int frame = 0; pathtrace(frame, iteration); } if (ui_showGbuffer) { - showGBuffer(pbo_dptr); - } else { - showImage(pbo_dptr, iteration); + showGBuffer(pbo_dptr, ui_currGBuffer); + } + else if (ui_denoise) { + denoise(pbo_dptr, iteration, ui_colorWeight, ui_normalWeight, ui_positionWeight, ui_filterSize, ui_atrous, denoiserCallCount); + denoiserCallCount++; + } + else { + showImage(pbo_dptr, iteration); + ptCallCount++; + } + timer.endCpuTimer(); + + // for performance analysis + if (iteration == ui_iterations) { + if (ui_denoise && denoiserCallCount < TIMER_COUNT) { + std::cout << "Path tracer with denoiser execution: " << timer.getCpuElapsedTimeForPreviousOperation() << "ms." << std::endl; + } + else if (ptCallCount < TIMER_COUNT) { + std::cout << "Path tracer execution: " << timer.getCpuElapsedTimeForPreviousOperation() << "ms." << std::endl; + } } // unmap buffer object diff --git a/src/main.h b/src/main.h index 06d311a..e97f3bb 100644 --- a/src/main.h +++ b/src/main.h @@ -35,7 +35,9 @@ extern int height; extern int ui_iterations; extern int startupIterations; extern bool ui_showGbuffer; +extern int ui_currGBuffer; extern bool ui_denoise; +extern bool ui_atrous; extern int ui_filterSize; extern float ui_colorWeight; extern float ui_normalWeight; diff --git a/src/pathtrace.cu b/src/pathtrace.cu index 23e5f90..d087f21 100644 --- a/src/pathtrace.cu +++ b/src/pathtrace.cu @@ -16,6 +16,15 @@ #define ERRORCHECK 1 +// http://demofox.org/gauss.html with sigma=1.0, support=0.5 +const float kernel[25] = { + 0.003765, 0.015019, 0.023792, 0.015019, 0.003765, + 0.015019, 0.059912, 0.094907, 0.059912, 0.015019, + 0.023792, 0.094907, 0.150342, 0.094907, 0.023792, + 0.015019, 0.059912, 0.094907, 0.059912, 0.015019, + 0.003765, 0.015019, 0.023792, 0.015019, 0.003765, +}; + #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) { @@ -38,6 +47,12 @@ void checkCUDAErrorFn(const char *msg, const char *file, int line) { #endif } +PerformanceTimer& timer() +{ + static PerformanceTimer timer; + return timer; +} + __host__ __device__ thrust::default_random_engine makeSeededRandomEngine(int iter, int index, int depth) { int h = utilhash((1 << 31) | (depth << 22) | iter) ^ utilhash(index); @@ -67,18 +82,42 @@ __global__ void sendImageToPBO(uchar4* pbo, glm::ivec2 resolution, } } -__global__ void gbufferToPBO(uchar4* pbo, glm::ivec2 resolution, GBufferPixel* gBuffer) { +__global__ void gbufferToPBO( + uchar4* pbo, + glm::ivec2 resolution, + GBufferPixel* gBuffer, + int currGBuffer +) { int x = (blockIdx.x * blockDim.x) + threadIdx.x; int y = (blockIdx.y * blockDim.y) + threadIdx.y; if (x < resolution.x && y < resolution.y) { int index = x + (y * resolution.x); - float timeToIntersect = gBuffer[index].t * 256.0; - pbo[index].w = 0; - pbo[index].x = timeToIntersect; - pbo[index].y = timeToIntersect; - pbo[index].z = timeToIntersect; + // normal map + if (currGBuffer == 0) { + glm::vec3 color = gBuffer[index].normal * 0.5f + 0.5f; + color.x = glm::clamp((int)(color.x * 255.0), 0, 255); // color only goes from 0 to 255!!! + color.y = glm::clamp((int)(color.y * 255.0), 0, 255); + color.z = glm::clamp((int)(color.z * 255.0), 0, 255); + + pbo[index].w = 0; + pbo[index].x = color.x; + pbo[index].y = color.y; + pbo[index].z = color.z; + } + // position map + else if (currGBuffer == 1) { + glm::vec3 color = 0.1f * (gBuffer[index].position); + color.x = glm::clamp((int)(color.x * 255.0), 0, 255); // color only goes from 0 to 255!!! + color.y = glm::clamp((int)(color.y * 255.0), 0, 255); + color.z = glm::clamp((int)(color.z * 255.0), 0, 255); + + pbo[index].w = 0; + pbo[index].x = color.x; + pbo[index].y = color.y; + pbo[index].z = color.z; + } } } @@ -89,8 +128,9 @@ static Material * dev_materials = NULL; static PathSegment * dev_paths = NULL; static ShadeableIntersection * dev_intersections = NULL; static GBufferPixel* dev_gBuffer = NULL; -// TODO: static variables for device memory, any extra info you need, etc -// ... +static glm::vec3* dev_denoised_img_in = NULL; +static glm::vec3* dev_denoised_img_out = NULL; +static float* dev_kernel = NULL; void pathtraceInit(Scene *scene) { hst_scene = scene; @@ -113,7 +153,16 @@ void pathtraceInit(Scene *scene) { cudaMalloc(&dev_gBuffer, pixelcount * sizeof(GBufferPixel)); - // TODO: initialize any extra device memeory you need + // to store intermediate results after every denoising iteration + cudaMalloc(&dev_denoised_img_in, pixelcount * sizeof(glm::vec3)); + cudaMemset(dev_denoised_img_in, 0, pixelcount * sizeof(glm::vec3)); + + cudaMalloc(&dev_denoised_img_out, pixelcount * sizeof(glm::vec3)); + cudaMemset(dev_denoised_img_out, 0, pixelcount * sizeof(glm::vec3)); + + // placeholder for gaussian kernel used in denoiser + cudaMalloc(&dev_kernel, 25 * sizeof(float)); + cudaMemcpy(dev_kernel, kernel, 25 * sizeof(float), cudaMemcpyHostToDevice); checkCUDAError("pathtraceInit"); } @@ -125,7 +174,9 @@ void pathtraceFree() { cudaFree(dev_materials); cudaFree(dev_intersections); cudaFree(dev_gBuffer); - // TODO: clean up any extra device memory you created + cudaFree(dev_denoised_img_in); + cudaFree(dev_denoised_img_out); + cudaFree(dev_kernel); checkCUDAError("pathtraceFree"); } @@ -149,6 +200,7 @@ __global__ void generateRayFromCamera(Camera cam, int iter, int traceDepth, Path segment.ray.origin = cam.position; segment.color = glm::vec3(1.0f, 1.0f, 1.0f); + segment.throughput = glm::vec3(1.f); segment.ray.direction = glm::normalize(cam.view - cam.right * cam.pixelLength.x * ((float)x - (float)cam.resolution.x * 0.5f) @@ -221,6 +273,7 @@ __global__ void computeIntersections( intersections[path_index].t = t_min; intersections[path_index].materialId = geoms[hit_geom_index].materialid; intersections[path_index].surfaceNormal = normal; + intersections[path_index].position = intersect_point; } } } @@ -238,6 +291,7 @@ __global__ void shadeSimpleMaterials ( { ShadeableIntersection intersection = shadeableIntersections[idx]; PathSegment segment = pathSegments[idx]; + if (segment.remainingBounces == 0) { return; } @@ -249,16 +303,17 @@ __global__ void shadeSimpleMaterials ( Material material = materials[intersection.materialId]; glm::vec3 materialColor = material.color; + segment.throughput = materialColor; // If the material indicates that the object was a light, "light" the ray if (material.emittance > 0.0f) { - segment.color *= (materialColor * material.emittance); + segment.color *= (materialColor * material.emittance);// *segment.throughput; segment.remainingBounces = 0; } else { - segment.color *= materialColor; + segment.color *= segment.throughput; glm::vec3 intersectPos = intersection.t * segment.ray.direction + segment.ray.origin; - scatterRay(segment, intersectPos, intersection.surfaceNormal, material, rng); + scatterRayMoreMats(segment, intersectPos, intersection.surfaceNormal, material, rng); } // If there was no intersection, color the ray black. // Lots of renderers use 4 channel color, RGBA, where A = alpha, often @@ -281,10 +336,84 @@ __global__ void generateGBuffer ( int idx = blockIdx.x * blockDim.x + threadIdx.x; if (idx < num_paths) { - gBuffer[idx].t = shadeableIntersections[idx].t; + //gBuffer[idx].t = shadeableIntersections[idx].t; + gBuffer[idx].normal = shadeableIntersections[idx].surfaceNormal; + gBuffer[idx].position = shadeableIntersections[idx].position; } } +// An edge-avoiding a-trous denoising algorithm based on https://jo.dreggn.org/home/2010_atrous.pdf +__global__ void aTrousFilter( + glm::ivec2 resolution, + glm::vec3* inImage, + glm::vec3* outImage, + GBufferPixel* gBuffer, + float* kernel, + float c_phi, float n_phi, float p_phi, + float stepwidth, bool avoidEdge +) +{ + int idx_x = (blockIdx.x * blockDim.x) + threadIdx.x; + int idx_y = (blockIdx.y * blockDim.y) + threadIdx.y; + + if (idx_x >= resolution.x || idx_y >= resolution.y) + { + return; + } + + glm::vec3 sum = glm::vec3(0.f); // accumulate color + float sumWeight = 0.f; + + // read GBuffer values for the current pixel + int idx = idx_y * resolution.x + idx_x; + glm::vec3 cval = inImage[idx]; + glm::vec3 nval = gBuffer[idx].normal; + glm::vec3 pval = gBuffer[idx].position; + + for (int i = -2; i <=2 ; i++) // default to using a 5x5 kernel + { + for (int j = -2; j <= 2; j++) + { + // find neighboring pixel at desired offset (clamp to image boundaries) + int xOffset = idx_x + j * stepwidth; + int yOffset = idx_y + i * stepwidth; + glm::ivec2 uv = glm::clamp(glm::ivec2(xOffset, yOffset), glm::ivec2(0), resolution); + + float weight = 1.f; + float h = kernel[(i + 2) * 5 + (j + 2)]; + + // color difference between current and neighboring pixel + int currIdx = uv.x + uv.y * resolution.x; + glm::vec3 ctemp = inImage[currIdx]; + glm::vec3 t = cval - ctemp; + float dist2 = dot(t, t); + float c_w = min(exp(-dist2 / c_phi), 1.f); + + if (avoidEdge) + { + // normal difference + glm::vec3 ntemp = gBuffer[currIdx].normal; + t = nval - ntemp; + dist2 = max(dot(t, t) / (stepwidth * stepwidth), 0.f); + float n_w = min(exp(-dist2 / n_phi), 1.f); + + // position difference + glm::vec3 ptemp = gBuffer[currIdx].position; + t = pval - ptemp; + dist2 = dot(t, t); + float p_w = min(exp(-dist2 / p_phi), 1.f); + + // calculate weights + weight = c_w * n_w * p_w; + } + + sum += ctemp * weight * h; + sumWeight += weight * h; + } + } + outImage[idx] = sum / sumWeight; +} + // Add the current iteration's output to the overall image __global__ void finalGather(int nPaths, glm::vec3 * image, PathSegment * iterationPaths) { @@ -410,7 +539,7 @@ void pathtrace(int frame, int iter) { } // CHECKITOUT: this kernel "post-processes" the gbuffer/gbuffers into something that you can visualize for debugging. -void showGBuffer(uchar4* pbo) { +void showGBuffer(uchar4* pbo, int currGBuffer) { const Camera &cam = hst_scene->state.camera; const dim3 blockSize2d(8, 8); const dim3 blocksPerGrid2d( @@ -418,7 +547,7 @@ void showGBuffer(uchar4* pbo) { (cam.resolution.y + blockSize2d.y - 1) / blockSize2d.y); // CHECKITOUT: process the gbuffer results and send them to OpenGL buffer for visualization - gbufferToPBO<<>>(pbo, cam.resolution, dev_gBuffer); + gbufferToPBO<<>>(pbo, cam.resolution, dev_gBuffer, currGBuffer); } void showImage(uchar4* pbo, int iter) { @@ -431,3 +560,43 @@ const Camera &cam = hst_scene->state.camera; // Send results to OpenGL buffer for rendering sendImageToPBO<<>>(pbo, cam.resolution, iter, dev_image); } + +// denoise and show image +void denoise(uchar4* pbo, int iter, float c_phi, float n_phi, float p_phi, float filterSize, bool avoidEdge, int callCount) { + 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); + + // calculate iterations based on filter size + int numIter = filterSize < 5 ? 0 : floor(log2(filterSize / 5.f)); + + timer().startGpuTimer(); + + // copy initial input image + int pixelCount = cam.resolution.x * cam.resolution.y; + cudaMemcpy(dev_denoised_img_in, dev_image, pixelCount * sizeof(glm::vec3), cudaMemcpyDeviceToDevice); + + for (int i = 0; i <= numIter; i++) { + float stepwidth = 1 << i; + aTrousFilter << > > (cam.resolution, dev_denoised_img_in, dev_denoised_img_out, + dev_gBuffer, dev_kernel, c_phi, n_phi, p_phi, stepwidth, avoidEdge); + + // ping pong buffer + if (i != numIter - 1) { // don't swap at the last iteration + std::swap(dev_denoised_img_in, dev_denoised_img_out); + } + } + timer().endGpuTimer(); + + // for screenshots + cudaMemcpy(hst_scene->state.image.data(), dev_denoised_img_out, pixelCount * sizeof(glm::vec3), cudaMemcpyDeviceToHost); + + // Send results to OpenGL buffer for rendering + sendImageToPBO << > > (pbo, cam.resolution, iter, dev_denoised_img_out); + + if (callCount < TIMER_COUNT) { + std::cout << "Denoiser: " << timer().getGpuElapsedTimeForPreviousOperation() << "ms. " << std::endl; + } +} diff --git a/src/pathtrace.h b/src/pathtrace.h index 9e12f44..f665517 100644 --- a/src/pathtrace.h +++ b/src/pathtrace.h @@ -6,5 +6,6 @@ void pathtraceInit(Scene *scene); void pathtraceFree(); void pathtrace(int frame, int iteration); -void showGBuffer(uchar4 *pbo); +void showGBuffer(uchar4 *pbo, int currGBuffer); void showImage(uchar4 *pbo, int iter); +void denoise(uchar4* pbo, int iter, float c_phi, float n_phi, float p_phi, float filterSize, bool avoidEdge, int callCount); \ No newline at end of file diff --git a/src/preview.cpp b/src/preview.cpp index 3ca2718..af434ca 100644 --- a/src/preview.cpp +++ b/src/preview.cpp @@ -213,15 +213,17 @@ void drawGui(int windowWidth, int windowHeight) { ImGui::SliderInt("Iterations", &ui_iterations, 1, startupIterations); ImGui::Checkbox("Denoise", &ui_denoise); + ImGui::Checkbox("Preserve Edges", &ui_atrous); - ImGui::SliderInt("Filter Size", &ui_filterSize, 0, 100); - ImGui::SliderFloat("Color Weight", &ui_colorWeight, 0.0f, 10.0f); + ImGui::SliderInt("Filter Size", &ui_filterSize, 0, 200); + ImGui::SliderFloat("Color Weight", &ui_colorWeight, 0.0f, 200.0f); ImGui::SliderFloat("Normal Weight", &ui_normalWeight, 0.0f, 10.0f); ImGui::SliderFloat("Position Weight", &ui_positionWeight, 0.0f, 10.0f); ImGui::Separator(); ImGui::Checkbox("Show GBuffer", &ui_showGbuffer); + ImGui::Combo("GBuffer", &ui_currGBuffer, "Normal\0Position\0"); ImGui::Separator(); @@ -238,6 +240,7 @@ void drawGui(int windowWidth, int windowHeight) { void mainLoop() { while (!glfwWindowShouldClose(window)) { glfwPollEvents(); + runCuda(); string title = "CIS565 Path Tracer | " + utilityCore::convertIntToString(iteration) + " Iterations"; diff --git a/src/sceneStructs.h b/src/sceneStructs.h index da7e558..0d27ee2 100644 --- a/src/sceneStructs.h +++ b/src/sceneStructs.h @@ -62,6 +62,7 @@ struct RenderState { struct PathSegment { Ray ray; glm::vec3 color; + glm::vec3 throughput; int pixelIndex; int remainingBounces; }; @@ -72,11 +73,14 @@ struct PathSegment { struct ShadeableIntersection { float t; glm::vec3 surfaceNormal; + glm::vec3 position; int materialId; }; // CHECKITOUT - a simple struct for storing scene geometry information per-pixel. // What information might be helpful for guiding a denoising filter? struct GBufferPixel { - float t; + //float t; + glm::vec3 normal; + glm::vec3 position; }; diff --git a/src/utilities.h b/src/utilities.h index abb4f27..a9e669c 100644 --- a/src/utilities.h +++ b/src/utilities.h @@ -9,10 +9,16 @@ #include #include +#include +#include +#include + #define PI 3.1415926535897932384626422832795028841971f +#define INV_PI 0.31830988618379067f #define TWO_PI 6.2831853071795864769252867665590057683943f #define SQRT_OF_ONE_THIRD 0.5773502691896257645091487805019574556476f #define EPSILON 0.00001f +#define TIMER_COUNT 10 namespace utilityCore { extern float clamp(float f, float min, float max); @@ -24,3 +30,96 @@ namespace utilityCore { extern std::string convertIntToString(int number); extern std::istream& safeGetline(std::istream& is, std::string& t); //Thanks to http://stackoverflow.com/a/6089413 } + +/** +* From project 2 scan & stream compaction +* This class is used for timing the performance +* Uncopyable and unmovable +* +* Adapted from WindyDarian(https://github.com/WindyDarian) +*/ +class PerformanceTimer +{ +public: + PerformanceTimer() + { + cudaEventCreate(&event_start); + cudaEventCreate(&event_end); + } + + ~PerformanceTimer() + { + cudaEventDestroy(event_start); + cudaEventDestroy(event_end); + } + + void startCpuTimer() + { + if (cpu_timer_started) { throw std::runtime_error("CPU timer already started"); } + cpu_timer_started = true; + + time_start_cpu = std::chrono::high_resolution_clock::now(); + } + + void endCpuTimer() + { + time_end_cpu = std::chrono::high_resolution_clock::now(); + + if (!cpu_timer_started) { throw std::runtime_error("CPU timer not started"); } + + std::chrono::duration duro = time_end_cpu - time_start_cpu; + prev_elapsed_time_cpu_milliseconds = + static_cast(duro.count()); + + cpu_timer_started = false; + } + + void startGpuTimer() + { + if (gpu_timer_started) { throw std::runtime_error("GPU timer already started"); } + gpu_timer_started = true; + + cudaEventRecord(event_start); + } + + void endGpuTimer() + { + cudaEventRecord(event_end); + cudaEventSynchronize(event_end); + + if (!gpu_timer_started) { throw std::runtime_error("GPU timer not started"); } + + cudaEventElapsedTime(&prev_elapsed_time_gpu_milliseconds, event_start, event_end); + gpu_timer_started = false; + } + + float getCpuElapsedTimeForPreviousOperation() //noexcept + { + return prev_elapsed_time_cpu_milliseconds; + } + + float getGpuElapsedTimeForPreviousOperation() //noexcept + { + return prev_elapsed_time_gpu_milliseconds; + } + + // remove copy and move functions + PerformanceTimer(const PerformanceTimer&) = delete; + PerformanceTimer(PerformanceTimer&&) = delete; + PerformanceTimer& operator=(const PerformanceTimer&) = delete; + PerformanceTimer& operator=(PerformanceTimer&&) = delete; + +private: + cudaEvent_t event_start = nullptr; + cudaEvent_t event_end = nullptr; + + using time_point_t = std::chrono::high_resolution_clock::time_point; + time_point_t time_start_cpu; + time_point_t time_end_cpu; + + bool cpu_timer_started = false; + bool gpu_timer_started = false; + + float prev_elapsed_time_cpu_milliseconds = 0.f; + float prev_elapsed_time_gpu_milliseconds = 0.f; +}; \ No newline at end of file