diff --git a/README.md b/README.md index 110697c..5d35fa5 100644 --- a/README.md +++ b/README.md @@ -1,13 +1,78 @@ CUDA Path Tracer ================ -**University of Pennsylvania, CIS 565: GPU Programming and Architecture, Project 3** +**University of Pennsylvania, CIS 565: GPU Programming and Architecture** -* (TODO) YOUR NAME HERE -* Tested on: (TODO) Windows 22, i7-2222 @ 2.22GHz 22GB, GTX 222 222MB (Moore 2222 Lab) +**Anantha Srinivas** +[LinkedIn](https://www.linkedin.com/in/anantha-srinivas-00198958/), [Twitter](https://twitter.com/an2tha) -### (TODO: Your README) +**Tested on:** +* Windows 10, i7-8700 @ 3.20GHz 16GB, GTX 1080 8097MB (Personal) +* Built for Visual Studio 2017 using the v140 toolkit +--- -*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. +![Figure 0: Cover](img/infinite_area.png) +![Figure 0: mats](img/Materials.png) + +Introduction +--- + +Path tracer is computer graphics technique using which virtual scenes are rendered. The most popular version called as Monte Carlo path tracer is similar to the traditional Ray tracer, but instead of one ray per pixel multiple ray are launched. At each point of intersection, the rays are bounced off in a random direction, based on the material the ray intersected with. The distribution of these bounces are defined by a functional called as Bi Direction scattering distribution function. + +A Path tracer implemented on the CPU is highly inefficient, given that for each pixel, we will need to cast multiple rays from the camera and trace it around the scene. A GPU based implementation would speed up this process, with computation of color for each pixel being run almost simulatanously on multiple threads. + +![Figure 1: Ray Intersection](img/path_trace.png) + +Implemented Features +--- + +1. A lot of structural changes to the code to reflect the architecture as described in **Physically Based Rendering, Second Edition: From Theory To Implementation** by Pharr, Matt and Humphreys, Greg. + +2. Naive Integrator. + +3. The path tracer supports the following materials: +* Diffuse materials + +![Figure 2: Diffuse](img/Cornell_diffuse.png) + +* Refractive material +![Figure 3: Refractive](img/Cornel_refraction_reflection.png) + +* Specular (Perfectly specular and glossy) +![Figure 4: Parallel Mirrors](img/Glass_ball.png) + +4. Rays between each iteration are compacted to remove any dead rays. This is done to improve the performance. + +5. The Path segments are sorted by material type. This ensure that almost all warps executed on the GPU have similar execution time. + +6. The first intersection is cached for subsequent iterations. This saves redundant calculation on the GPU. + +7. Implemented Refraction (the refractive index is assumed to be 1.52, which should be later taken as input from the scene file) + +8. Stocahstic sampled AntiAliasing. This jitters the ray direction from camera by a small amount. Over multiple interations, this results in a blurred effect. + + +|| No Anti-Aliasing | Anti-Alias 1 | Anti-Alias 2 | +:-------------------------:|:-------------------------:|:-------------------------: |:-------------------------: + Original Image|![](img/AA_00.png) | ![](img/AA_01.png) | ![](img/AA_02.png) +Zoomed|![](img/AA_00Zoomed.png) | ![](img/AA_01Zoomed.png) | ![](img/AA_02Zoomed.png) + + +9. Implemented Procedural geometry or implicit surfaces. Current supported implicit surfaces include - box, sphere, torus, capped cylinder. The intersection with the geometry is determined using ray marching. + +![Figure 5: Implicit Surface](img/Implicit.png) + + +Analysis +--- +Comparison in runtimes between CPU vs GPU based Path tracer (for Naive Integrator) + +When compared to GPU based Path tracer, CPU is way more slow. Even with multi-threading enabled, the number of pixels computed per thread is limited. + +![Figure 6: Performance grah](img/performance.png) + + +References +--- +[PBRT] Physically Based Rendering, Second Edition: From Theory To Implementation. Pharr, Matt and Humphreys, Greg. 2010. \ No newline at end of file diff --git a/img/AA_00.png b/img/AA_00.png new file mode 100644 index 0000000..66024df Binary files /dev/null and b/img/AA_00.png differ diff --git a/img/AA_00Zoomed.png b/img/AA_00Zoomed.png new file mode 100644 index 0000000..8a47ae0 Binary files /dev/null and b/img/AA_00Zoomed.png differ diff --git a/img/AA_01.png b/img/AA_01.png new file mode 100644 index 0000000..c4829f0 Binary files /dev/null and b/img/AA_01.png differ diff --git a/img/AA_01Zoomed.png b/img/AA_01Zoomed.png new file mode 100644 index 0000000..a712a92 Binary files /dev/null and b/img/AA_01Zoomed.png differ diff --git a/img/AA_02.png b/img/AA_02.png new file mode 100644 index 0000000..b0ce13a Binary files /dev/null and b/img/AA_02.png differ diff --git a/img/AA_02Zoomed.png b/img/AA_02Zoomed.png new file mode 100644 index 0000000..d280469 Binary files /dev/null and b/img/AA_02Zoomed.png differ diff --git a/img/Cornel_refraction_reflection.png b/img/Cornel_refraction_reflection.png new file mode 100644 index 0000000..2b07c3f Binary files /dev/null and b/img/Cornel_refraction_reflection.png differ diff --git a/img/Cornell_diffuse.png b/img/Cornell_diffuse.png new file mode 100644 index 0000000..79dd51e Binary files /dev/null and b/img/Cornell_diffuse.png differ diff --git a/img/Glass_ball.png b/img/Glass_ball.png new file mode 100644 index 0000000..fe8edcc Binary files /dev/null and b/img/Glass_ball.png differ diff --git a/img/Implicit.png b/img/Implicit.png new file mode 100644 index 0000000..64456c8 Binary files /dev/null and b/img/Implicit.png differ diff --git a/img/Materials.png b/img/Materials.png new file mode 100644 index 0000000..a3b6863 Binary files /dev/null and b/img/Materials.png differ diff --git a/img/Mirror_room.png b/img/Mirror_room.png new file mode 100644 index 0000000..fc47efd Binary files /dev/null and b/img/Mirror_room.png differ diff --git a/img/cornell_iteration_100_depth_8_camera_800x800.png b/img/cornell_iteration_100_depth_8_camera_800x800.png new file mode 100644 index 0000000..47886b1 Binary files /dev/null and b/img/cornell_iteration_100_depth_8_camera_800x800.png differ diff --git a/img/cornell_iteration_100_depth_8_camera_800x800_reflective.png b/img/cornell_iteration_100_depth_8_camera_800x800_reflective.png new file mode 100644 index 0000000..f88f276 Binary files /dev/null and b/img/cornell_iteration_100_depth_8_camera_800x800_reflective.png differ diff --git a/img/infinite_2.png b/img/infinite_2.png new file mode 100644 index 0000000..ff988cb Binary files /dev/null and b/img/infinite_2.png differ diff --git a/img/infinite_area.png b/img/infinite_area.png new file mode 100644 index 0000000..a84909f Binary files /dev/null and b/img/infinite_area.png differ diff --git a/img/path_trace.png b/img/path_trace.png new file mode 100644 index 0000000..10e2a12 Binary files /dev/null and b/img/path_trace.png differ diff --git a/img/performance.png b/img/performance.png new file mode 100644 index 0000000..3a2d287 Binary files /dev/null and b/img/performance.png differ diff --git a/scenes/infinite_room.txt b/scenes/infinite_room.txt new file mode 100644 index 0000000..abe6454 --- /dev/null +++ b/scenes/infinite_room.txt @@ -0,0 +1,144 @@ +// 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 + +// Specular white mirror +MATERIAL 5 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB 1.0 1.0 1.0 +REFL 1 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Refractive glass ball +MATERIAL 6 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB 1.0 1.0 1.0 +REFL 0 +REFR 1 +REFRIOR 0 +EMITTANCE 0 + +// Camera +CAMERA +RES 800 800 +FOVY 45 +ITERATIONS 500 +DEPTH 10 +FILE cornell +EYE -0.5 4 2 +LOOKAT -10 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 5 +TRANS 0 5 -5 +ROTAT 0 90 0 +SCALE .01 10 10 + +// Left wall +OBJECT 4 +cube +material 5 +TRANS -5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +// Right wall +OBJECT 5 +plane +material 5 +TRANS 5 5 0 +ROTAT 0 -90 0 +SCALE 10 10 1 + +// Front wall +OBJECT 6 +plane +material 5 +TRANS 0 5 5 +ROTAT 0 180 0 +SCALE 10 10 1 + +// Sphere +OBJECT 7 +sphere +material 2 +TRANS 0 4 -1 +ROTAT 0 0 0 +SCALE 3 3 3 diff --git a/scenes/test.txt b/scenes/test.txt new file mode 100644 index 0000000..820063b --- /dev/null +++ b/scenes/test.txt @@ -0,0 +1,166 @@ +// Emissive material (light) +MATERIAL 0 +RGB 1 1 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 5 +DIFFUSE 0 + +// Diffuse white +MATERIAL 1 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 +DIFFUSE 1 + +// Diffuse red +MATERIAL 2 +RGB .85 .35 .35 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 +DIFFUSE 1 + +// Diffuse green +MATERIAL 3 +RGB .35 .85 .35 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 +DIFFUSE 1 + +// Specular white +MATERIAL 4 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB .98 .98 .98 +REFL 1 +REFR 0 +REFRIOR 0 +EMITTANCE 0 +DIFFUSE 0 + +// Specular white mirror +MATERIAL 5 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB 1.0 1.0 1.0 +REFL 1 +REFR 0 +REFRIOR 0 +EMITTANCE 0 +DIFFUSE 0 + +// Refractive glass ball +MATERIAL 6 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB 0.98 0.98 0.98 +REFL 0 +REFR 1 +REFRIOR 0 +EMITTANCE 0 +DIFFUSE 0 + +// Diffuse yellow +MATERIAL 7 +RGB 1.0 .85 .32 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 +DIFFUSE 1 + +// Plastic blue +MATERIAL 8 +RGB 0.08 0.15 0.6 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 +DIFFUSE 1 + +// Camera +CAMERA +RES 800 800 +FOVY 45 +ITERATIONS 100 +DEPTH 10 +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 +plane +material 3 +TRANS 5 5 0 +ROTAT 0 -90 0 +SCALE 10 10 1 + +// Sphere diffuse +OBJECT 6 +cube +material 7 +TRANS 0 4 -1 +ROTAT 45 45 0 +SCALE 3 3 3 + diff --git a/scenes/test2.txt b/scenes/test2.txt new file mode 100644 index 0000000..691d4e4 --- /dev/null +++ b/scenes/test2.txt @@ -0,0 +1,182 @@ +// Emissive material (light) +MATERIAL 0 +RGB 1 1 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 5 +DIFFUSE 0 + +// Diffuse white +MATERIAL 1 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 +DIFFUSE 1 + +// Diffuse red +MATERIAL 2 +RGB .85 .35 .35 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 +DIFFUSE 1 + +// Diffuse green +MATERIAL 3 +RGB .35 .85 .35 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 +DIFFUSE 1 + +// Specular white +MATERIAL 4 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB .98 .98 .98 +REFL 1 +REFR 0 +REFRIOR 0 +EMITTANCE 0 +DIFFUSE 0 + +// Specular white mirror +MATERIAL 5 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB 1.0 1.0 1.0 +REFL 1 +REFR 0 +REFRIOR 0 +EMITTANCE 0 +DIFFUSE 0 + +// Refractive glass ball +MATERIAL 6 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB 1.0 1.0 1.0 +REFL 0 +REFR 1 +REFRIOR 0 +EMITTANCE 0 +DIFFUSE 0 + +// Diffuse yellow +MATERIAL 7 +RGB 1.0 .85 .32 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 +DIFFUSE 1 + +// Plastic blue +MATERIAL 8 +RGB 0.08 0.15 0.6 +SPECEX 0 +SPECRGB 0 0 0 +REFL 1 +REFR 0 +REFRIOR 0 +EMITTANCE 0 +DIFFUSE 1 + +// Camera +CAMERA +RES 800 800 +FOVY 45 +ITERATIONS 1000 +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 +plane +material 3 +TRANS 5 5 0 +ROTAT 0 -90 0 +SCALE 10 10 1 + +// Sphere mirror +OBJECT 6 +sphere +material 5 +TRANS -2 2 -1 +ROTAT 0 0 0 +SCALE 2 2 2 + +// Sphere diffuse +OBJECT 7 +sphere +material 7 +TRANS 0 2 -1 +ROTAT 0 0 0 +SCALE 2 2 2 + +// Sphere plastic +OBJECT 8 +sphere +material 8 +TRANS 2 2 -1 +ROTAT 0 0 0 +SCALE 2 2 2 + diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index a1cb3fb..84d8a59 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -19,5 +19,5 @@ set(SOURCE_FILES cuda_add_library(src ${SOURCE_FILES} - OPTIONS -arch=sm_20 + OPTIONS -arch=sm_60 ) diff --git a/src/bsdf.cu b/src/bsdf.cu new file mode 100644 index 0000000..223a8dc --- /dev/null +++ b/src/bsdf.cu @@ -0,0 +1,213 @@ +#include +#include "glm/glm.hpp" + +#include "sceneStructs.h" +#include "warpfunctions.cu" + + + +namespace BSDF +{ + namespace + { + // Lamberts BRDF ------------------------------------------------------- + __host__ __device__ glm::vec3 Lamberts_F(const glm::vec3* wo, const glm::vec3* wi, const Material* material) + { + return material->color * InvPi; + } + + __host__ __device__ float Lamberts_Pdf(const glm::vec3* wo, glm::vec3* wi) + { + return Common::SameHemisphere(wo, wi) ? Common::AbsCosTheta(wi) * InvPi : 0; + } + + __host__ __device__ glm::vec3 Lamberts_SampleF(const glm::vec3* wo, glm::vec3* wi, float* pdf, const glm::vec2* xi, const Material* material) + { + // 1. Cosine sample the hemisphere + *wi = WarpFunctions::SquareToHemisphereCosine(xi); + + if (wo->z < 0) { + wi->z *= -1; + } + + // 2. Calculate the pdf + *pdf = Lamberts_Pdf(wo, wi); + + // 3. return f + return Lamberts_F(wo, wi, material); + } + + // Specular BRDF ------------------------------------------------------- + __host__ __device__ glm::vec3 SpecularR_F(const glm::vec3* wo, const glm::vec3* wi, const Material* material) + { + return material->color; + } + + __host__ __device__ float SpecularR_Pdf(const glm::vec3* wo, glm::vec3* wi) + { + return 1.f; + } + + __host__ __device__ glm::vec3 SpecularR_SampleF(const glm::vec3* wo, glm::vec3* wi, float* pdf, const glm::vec2* xi, const Material* material) + { + *wi = glm::vec3(-(*wo).x, -(*wo).y, (*wo).z); + *pdf = SpecularR_Pdf(wo, wi); + return SpecularR_F(wo, wi, material); + } + + // Specular BTDF ------------------------------------------------------- + + __host__ __device__ glm::vec3 SpecularT_F(const glm::vec3* wo, const glm::vec3* wi, const Material* material) + { + return material->color; + } + + __host__ __device__ float SpecularT_Pdf(const glm::vec3* wo, glm::vec3* wi) + { + return 1.f; + } + + __host__ __device__ glm::vec3 SpecularT_SampleF(const glm::vec3* wo, glm::vec3* wi, float* pdf, const glm::vec2* xi, const Material* material) + { + const bool entering = wo->z > 0; + + const float etaA = material->refractive.etaA; + const float etaB = material->refractive.etaB; + + const float etaI = entering ? etaA : etaB; + const float etaT = entering ? etaB : etaA; + + // Check if refraction has occured (to see for total internal reflection) + const bool refracted = Common::Refract(*wo, Common::Faceforward(glm::vec3(0.f, 0.f, 1.f), *wo), etaI / etaT, wi); + + if (!refracted) + return glm::vec3(0.f); + + *pdf = 1.f; + + // TODO :: Add Fresnel + const glm::vec3 ft = material->color / Common::AbsCosTheta(wi); + + return ft; + } + + // TODO: Add more bxdfs + + } // Anonymous namespace end + + __host__ __device__ glm::vec3 Sample_F(glm::vec3 woW, glm::vec3* wiW, float* pdf, glm::vec2* xi, const Material* material, ShadeableIntersection* intersection) + { + // TODO: + + // 1. Select Random Bxdf + const int numBxDFs = material->numBxdfs; + // TODO : This can be done later as we are using only one material + + int randomBxdf = int((*xi)[0] * numBxDFs) % numBxDFs; + BxDFType selBxdf = material->bxdfTypes[randomBxdf]; + + // 2. Rewriting the random number + glm::vec2 temp = glm::vec2((*xi)[0], (*xi)[1]); + + // 3. Converting wo, wi to tangent space + const glm::vec3 woL = intersection->m_worldToTangent * (woW); + glm::vec3 wiL;// = worldToTangent * (*wiW); + + // 4. Getting the color of the random bxdf + + glm::vec3 selBxdfCol(0.f); + + // TODO : Optimize this + if(selBxdf == BxDFType::BSDF_REFLECTION) + { + selBxdfCol = SpecularR_SampleF(&woL, &wiL, pdf, &temp, material); + } + else if(selBxdf == BxDFType::BSDF_TRANSMISSION) + { + selBxdfCol = SpecularT_SampleF(&woL, &wiL, pdf, &temp, material); + } + else if(selBxdf == BxDFType::BSDF_DIFFUSE) + { + selBxdfCol = Lamberts_SampleF(&woL, &wiL, pdf, &temp, material); + } + + const glm::vec3 wow = intersection->m_tangentToWorld * wiL; + + *wiW = wow; + + if(selBxdf == BxDFType::BSDF_REFLECTION || selBxdf == BxDFType::BSDF_TRANSMISSION) + { + return selBxdfCol; + } + + return selBxdfCol; + } + + __host__ __device__ glm::vec3 F(const glm::vec3* woW, const glm::vec3* wiW, const Material* material, const ShadeableIntersection* intersection, const BxDFType flags) + { + glm::vec3 color(0.f); + + glm::vec3 woL = intersection->m_worldToTangent * (*woW); + glm::vec3 wiL = intersection->m_worldToTangent * (*wiW); + + /*for (int i = 0; i < numBxDFs; ++i) { + sum += bxdfs[i]->MatchesFlags(flags) ? bxdfs[i]->f(woL, wiL) : Color3f(0.f); + }*/ + + // TODO : This can be done later as we are using only one material + if(material->hasReflective) + { + color += SpecularR_F(&woL, &wiL, material); + } + else if(material->hasRefractive) + { + color += SpecularT_F(&woL, &wiL, material); + } + else + { + color += Lamberts_F(&woL, &wiL, material); + } + + return color; + } + + __host__ __device__ float Pdf(const glm::vec3* woW, const glm::vec3* wiW, const Material* material, const ShadeableIntersection* intersection, const BxDFType flags) + { + int numPdfs = 0; + float sumPdf = 0; + + // Converting them to tangent space before sending to bxdf pdf + glm::vec3 woL = intersection->m_worldToTangent * (*woW); + glm::vec3 wiL = intersection->m_worldToTangent * (*wiW); + + /*for (int i = 0; i < numBxDFs; ++i) { + if (bxdfs[i]->MatchesFlags(flags)) { + sumPdf += bxdfs[i]->Pdf(woL, wiL); + numPdfs++; + } + }*/ + // TODO : This can be done later as we are using only one material + if(material->hasReflective) + { + sumPdf = SpecularR_Pdf(&woL, &wiL); + } + else if(material->hasRefractive) + { + sumPdf = SpecularT_Pdf(&woL, &wiL); + } + else + { + sumPdf = Lamberts_Pdf(&woL, &wiL); + } + + //sumPdf /= (1.f * numPdfs); + return sumPdf; + } +} // namespace BSDF end + + + + + + + diff --git a/src/bxdf.cu b/src/bxdf.cu new file mode 100644 index 0000000..ac170ae --- /dev/null +++ b/src/bxdf.cu @@ -0,0 +1,111 @@ +#pragma once + +#include +#include "glm/glm.hpp" + +#include "common.h" +#include "warpfunctions.h" + +enum BxDFType +{ + BSDF_REFLECTION = 1 << 0, // This BxDF handles rays that are reflected off surfaces + BSDF_TRANSMISSION = 1 << 1, // This BxDF handles rays that are transmitted through surfaces + BSDF_DIFFUSE = 1 << 2, // This BxDF represents diffuse energy scattering, which is uniformly random + BSDF_GLOSSY = 1 << 3, // This BxDF represents glossy energy scattering, which is biased toward certain directions + BSDF_SPECULAR = 1 << 4, // This BxDF handles specular energy scattering, which has no element of randomness + BSDF_ALL = BSDF_DIFFUSE | BSDF_GLOSSY | BSDF_SPECULAR | BSDF_REFLECTION | BSDF_TRANSMISSION +}; + +//--------------------------------------------------------------------------------- +// BxDF +//--------------------------------------------------------------------------------- +struct Bxdf +{ + int id = 0; + + Bxdf() {} + + __device__ virtual glm::vec3 SampleF(const glm::vec3* wo, glm::vec3* wi, float* pdf, const glm::vec2* xi, const glm::vec3 color) + { + return glm::vec3(1.f); + } + + __device__ virtual float Pdf(const glm::vec3* wo, glm::vec3* wi) + { + return 0.f; + } + + __device__ virtual glm::vec3 F(const glm::vec3* wo, const glm::vec3* wi, const glm::vec3 color) + { + return glm::vec3(0.f); + } +}; + +//--------------------------------------------------------------------------------- +// Lamberts or Diffuse BRDF +//--------------------------------------------------------------------------------- +struct LambertsBrdf : public Bxdf +{ + + LambertsBrdf() + { + id = 1; + } + + // Lamberts BRDF ------------------------------------------------------- + __device__ glm::vec3 F(const glm::vec3* wo, const glm::vec3* wi, const glm::vec3 color) + { + return color * InvPi; + } + + __device__ float Pdf(const glm::vec3* wo, glm::vec3* wi) + { + return Common::SameHemisphere(wo, wi) ? Common::AbsCosTheta(wi) * InvPi : 0; + } + + __device__ glm::vec3 SampleF(const glm::vec3* wo, glm::vec3* wi, float* pdf, const glm::vec2* xi, const glm::vec3 color) + { + // 1. Cosine sample the hemisphere + *wi = WarpFunctions::SquareToHemisphereCosine(xi); + + if (wo->z < 0) { + wi->z *= -1; + } + + // 2. Calculate the pdf + *pdf = Pdf(wo, wi); + + // 3. return f + return F(wo, wi, color); + } +}; + +//--------------------------------------------------------------------------------- +// Specular BRDF +//--------------------------------------------------------------------------------- +struct SpecularBrdf : public Bxdf +{ + + SpecularBrdf() + { + id = 2; + } + + // Specular BRDF ------------------------------------------------------- + __device__ glm::vec3 F(const glm::vec3* wo, const glm::vec3* wi, const glm::vec3 color) + { + return color; + } + + __device__ float Pdf(const glm::vec3* wo, glm::vec3* wi) + { + return 1.f; + } + + __device__ glm::vec3 SampleF(const glm::vec3* wo, glm::vec3* wi, float* pdf, const glm::vec2* xi, const glm::vec3 color) + { + *wi = glm::vec3(-(*wo).x, -(*wo).y, (*wo).z); + *pdf = Pdf(wo, wi); + return F(wo, wi, color); + } +}; diff --git a/src/common.cu b/src/common.cu new file mode 100644 index 0000000..0f3714e --- /dev/null +++ b/src/common.cu @@ -0,0 +1,51 @@ +#include +#include "glm/glm.hpp" + +namespace Common +{ + #define Pi 3.14159265358979323846f + #define TwoPi 6.28318530717958647692f + #define PiOver4 0.78539816339744830961f + + #define InvPi 0.31830988618379067154f + #define Inv2Pi 0.15915494309189533577f + #define Inv4Pi 0.07957747154594766788f + #define Inv8Pi 0.03978873577f + + __host__ __device__ inline float AbsCosTheta(const glm::vec3* w) + { + return glm::abs(w->z); + } + + __host__ __device__ inline bool SameHemisphere(const glm::vec3* w, const glm::vec3* wp) + { + return (w->z * wp->z > 0); + } + + __host__ __device__ inline glm::vec3 Faceforward(const glm::vec3 &n, const glm::vec3 &v) { + return (glm::dot(n, v) < 0.f) ? -n : n; + } + + + __host__ __device__ inline bool Refract(glm::vec3 wi, glm::vec3 n, float eta, glm::vec3 *wt) + { + // Compute cos theta using Snell's law + float cosThetaI = glm::dot(n, wi); + float sin2ThetaI = glm::max(float(0), float(1 - cosThetaI * cosThetaI)); + float sin2ThetaT = eta * eta * sin2ThetaI; + + // Handle total internal reflection for transmission + if (sin2ThetaT >= 1) return false; + float cosThetaT = glm::sqrt(1 - sin2ThetaT); + *wt = eta * -wi + (eta * cosThetaI - cosThetaT) * glm::vec3(n); + return true; + } + +} // namespace Common end + + + + + + + diff --git a/src/common.h b/src/common.h new file mode 100644 index 0000000..931aea2 --- /dev/null +++ b/src/common.h @@ -0,0 +1,20 @@ +#pragma once + +#include +#include "glm/glm.hpp" + +namespace Common +{ +#define Pi 3.14159265358979323846f +#define TwoPi 6.28318530717958647692f +#define PiOver4 0.78539816339744830961f + +#define InvPi 0.31830988618379067154f +#define Inv2Pi 0.15915494309189533577f +#define Inv4Pi 0.07957747154594766788f +#define Inv8Pi 0.03978873577f + + float AbsCosTheta(const glm::vec3* w); + bool SameHemisphere(const glm::vec3* w, const glm::vec3* wp); + +} // namespace Common end diff --git a/src/intersections.h b/src/intersections.h index 6f23872..45dcbb0 100644 --- a/src/intersections.h +++ b/src/intersections.h @@ -1,144 +1,3 @@ #pragma once -#include -#include -#include "sceneStructs.h" -#include "utilities.h" - -/** - * Handy-dandy hash function that provides seeds for random number generation. - */ -__host__ __device__ inline unsigned int utilhash(unsigned int a) { - a = (a + 0x7ed55d16) + (a << 12); - a = (a ^ 0xc761c23c) ^ (a >> 19); - a = (a + 0x165667b1) + (a << 5); - a = (a + 0xd3a2646c) ^ (a << 9); - a = (a + 0xfd7046c5) + (a << 3); - a = (a ^ 0xb55a4f09) ^ (a >> 16); - return a; -} - -// CHECKITOUT -/** - * Compute a point at parameter value `t` on ray `r`. - * Falls slightly short so that it doesn't intersect the object it's hitting. - */ -__host__ __device__ glm::vec3 getPointOnRay(Ray r, float t) { - return r.origin + (t - .0001f) * glm::normalize(r.direction); -} - -/** - * Multiplies a mat4 and a vec4 and returns a vec3 clipped from the vec4. - */ -__host__ __device__ glm::vec3 multiplyMV(glm::mat4 m, glm::vec4 v) { - return glm::vec3(m * v); -} - -// CHECKITOUT -/** - * Test intersection between a ray and a transformed cube. Untransformed, - * the cube ranges from -0.5 to 0.5 in each axis and is centered at the origin. - * - * @param intersectionPoint Output parameter for point of intersection. - * @param normal Output parameter for surface normal. - * @param outside Output param for whether the ray came from outside. - * @return Ray parameter `t` value. -1 if no intersection. - */ -__host__ __device__ float boxIntersectionTest(Geom box, Ray r, - glm::vec3 &intersectionPoint, glm::vec3 &normal, bool &outside) { - Ray q; - q.origin = multiplyMV(box.inverseTransform, glm::vec4(r.origin , 1.0f)); - q.direction = glm::normalize(multiplyMV(box.inverseTransform, glm::vec4(r.direction, 0.0f))); - - float tmin = -1e38f; - float tmax = 1e38f; - glm::vec3 tmin_n; - glm::vec3 tmax_n; - for (int xyz = 0; xyz < 3; ++xyz) { - float qdxyz = q.direction[xyz]; - /*if (glm::abs(qdxyz) > 0.00001f)*/ { - float t1 = (-0.5f - q.origin[xyz]) / qdxyz; - float t2 = (+0.5f - q.origin[xyz]) / qdxyz; - float ta = glm::min(t1, t2); - float tb = glm::max(t1, t2); - glm::vec3 n; - n[xyz] = t2 < t1 ? +1 : -1; - if (ta > 0 && ta > tmin) { - tmin = ta; - tmin_n = n; - } - if (tb < tmax) { - tmax = tb; - tmax_n = n; - } - } - } - - if (tmax >= tmin && tmax > 0) { - outside = true; - if (tmin <= 0) { - tmin = tmax; - tmin_n = tmax_n; - outside = false; - } - intersectionPoint = multiplyMV(box.transform, glm::vec4(getPointOnRay(q, tmin), 1.0f)); - normal = glm::normalize(multiplyMV(box.transform, glm::vec4(tmin_n, 0.0f))); - return glm::length(r.origin - intersectionPoint); - } - return -1; -} - -// CHECKITOUT -/** - * Test intersection between a ray and a transformed sphere. Untransformed, - * the sphere always has radius 0.5 and is centered at the origin. - * - * @param intersectionPoint Output parameter for point of intersection. - * @param normal Output parameter for surface normal. - * @param outside Output param for whether the ray came from outside. - * @return Ray parameter `t` value. -1 if no intersection. - */ -__host__ __device__ float sphereIntersectionTest(Geom sphere, Ray r, - glm::vec3 &intersectionPoint, glm::vec3 &normal, bool &outside) { - float radius = .5; - - glm::vec3 ro = multiplyMV(sphere.inverseTransform, glm::vec4(r.origin, 1.0f)); - glm::vec3 rd = glm::normalize(multiplyMV(sphere.inverseTransform, glm::vec4(r.direction, 0.0f))); - - Ray rt; - rt.origin = ro; - rt.direction = rd; - - float vDotDirection = glm::dot(rt.origin, rt.direction); - float radicand = vDotDirection * vDotDirection - (glm::dot(rt.origin, rt.origin) - powf(radius, 2)); - if (radicand < 0) { - return -1; - } - - float squareRoot = sqrt(radicand); - float firstTerm = -vDotDirection; - float t1 = firstTerm + squareRoot; - float t2 = firstTerm - squareRoot; - - float t = 0; - if (t1 < 0 && t2 < 0) { - return -1; - } else if (t1 > 0 && t2 > 0) { - t = min(t1, t2); - outside = true; - } else { - t = max(t1, t2); - outside = false; - } - - glm::vec3 objspaceIntersection = getPointOnRay(rt, t); - - intersectionPoint = multiplyMV(sphere.transform, glm::vec4(objspaceIntersection, 1.f)); - normal = glm::normalize(multiplyMV(sphere.invTranspose, glm::vec4(objspaceIntersection, 0.f))); - if (!outside) { - normal = -normal; - } - - return glm::length(r.origin - intersectionPoint); -} diff --git a/src/lights.cu b/src/lights.cu new file mode 100644 index 0000000..7884251 --- /dev/null +++ b/src/lights.cu @@ -0,0 +1,105 @@ +#pragma once + +#include +#include "glm/glm.hpp" +#include "glm/gtx/norm.hpp" + +#include "sceneStructs.h" +#include "shapes.cu" + +namespace Lights +{ + +#define EPSILON 0.001 +#define ISZERO(p) abs(p) < EPSILON + + __host__ __device__ bool isSamplePoint(const glm::vec3* p1, const glm::vec3* p2) + { + return glm::distance(*p1, *p2) < EPSILON; + } + + // Could potentially keep other type of lights + namespace DiffuseAreaLight + { + __host__ __device__ glm::vec3 L(const Material* material, const glm::vec3* point, const glm::vec3* normal) + { + // This seems to be giving the intersection with the light geometry + const bool isFront = glm::dot(*normal, -(*point)) > 0; + return isFront? (material->emittance * material->color) : glm::vec3(0.f); + } + + __host__ __device__ glm::vec3 Sample_Li(Geom* lightGeometry, glm::vec2* xi, glm::vec3* wiW, float* pdf, const Material* material, ShadeableIntersection* refIntersection) + { + + // 1. Get an Intersection on the surface of its Shape by invoking shape->Sample. + + // Get the intersection with the geometry of the object + glm::vec3 intersection_point(0.f); + glm::vec3 intersection_normal(0.f); + + // Currently this is the only implemented shape + if(lightGeometry->type == PLANE) + { + Shapes::SquarePlane::Sample(lightGeometry, xi, pdf, &intersection_point, &intersection_normal); + } + + if(isSamplePoint(&intersection_point, &refIntersection->m_intersectionPointWorld)) + { + *pdf = 0; + return glm::vec3(0.f); + } + + // Vector from reference point to out light's intersection point + const glm::vec3 wi = glm::normalize(intersection_point - intersection_point); + + // distance betweern light and intersection point + const float r = glm::distance2(intersection_point, refIntersection->m_intersectionPointWorld); + + // grazing angle + const float cosAngle = glm::abs(glm::dot(intersection_normal, -wi)); + + // initial area of shape of light + const float area = *pdf; + + // Can't divide by zero + if(cosAngle == 0.f) + { + *pdf = 0.f; + } + else + { + *pdf = (r) / (cosAngle / area); + } + + // 2. Check if the resultant PDF is zero or that the reference Intersection + // and the resultant Intersection are the same point in space, and return black if this is the case. + if(ISZERO(*pdf)) + { + return glm::vec3(0.f); + } + + // 3. Set wi to the normalized vector from the reference Intersection's point to the Shape's intersection point. + *wiW = glm::normalize(intersection_point - refIntersection->m_intersectionPointWorld); + + // 4. Return the light emitted along ωi from our intersection point + return L(material, wiW, &intersection_normal); + + } + + __host__ __device__ float Pdf_Li(ShadeableIntersection* intersection, const glm::vec3 &wi) + { + // TODO: + return 0.f; + //return shape->Pdf(ref, wi); + } + + } // namespace DiffuseAreaLight end + +} // namespace Lights end + + + + + + + diff --git a/src/main.cpp b/src/main.cpp index fe8e85e..2f87d04 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -35,6 +35,7 @@ int main(int argc, char** argv) { if (argc < 2) { printf("Usage: %s SCENEFILE.txt\n", argv[0]); + getchar(); return 1; } @@ -134,11 +135,12 @@ void runCuda() { // execute the kernel int frame = 0; - pathtrace(pbo_dptr, frame, iteration); + pathtrace(pbo_dptr, frame, iteration, renderState->iterations); // unmap buffer object cudaGLUnmapBufferObject(pbo); } else { + getchar(); saveImage(); pathtraceFree(); cudaDeviceReset(); diff --git a/src/pathtrace.cu b/src/pathtrace.cu index c1ec122..8cd2f71 100644 --- a/src/pathtrace.cu +++ b/src/pathtrace.cu @@ -4,20 +4,25 @@ #include #include #include +#include +#include #include "sceneStructs.h" #include "scene.h" #include "glm/glm.hpp" -#include "glm/gtx/norm.hpp" #include "utilities.h" #include "pathtrace.h" #include "intersections.h" #include "interactions.h" +#include "bsdf.cu" +#include "shapes.cu" #define ERRORCHECK 1 #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(); @@ -38,6 +43,19 @@ void checkCUDAErrorFn(const char *msg, const char *file, int line) { #endif } +/** +* Handy-dandy hash function that provides seeds for random number generation. +*/ +__host__ __device__ inline unsigned int utilhash(unsigned int a) { + a = (a + 0x7ed55d16) + (a << 12); + a = (a ^ 0xc761c23c) ^ (a >> 19); + a = (a + 0x165667b1) + (a << 5); + a = (a + 0xd3a2646c) ^ (a << 9); + a = (a + 0xfd7046c5) + (a << 3); + a = (a ^ 0xb55a4f09) ^ (a >> 16); + return a; +} + __host__ __device__ thrust::default_random_engine makeSeededRandomEngine(int iter, int index, int depth) { int h = utilhash((1 << 31) | (depth << 22) | iter) ^ utilhash(index); @@ -67,15 +85,51 @@ __global__ void sendImageToPBO(uchar4* pbo, glm::ivec2 resolution, } } +#define PDF_EPSILON 0.001 +#define RAY_EPSILON 0.0005f + static Scene * hst_scene = NULL; static glm::vec3 * dev_image = NULL; static Geom * dev_geoms = NULL; static Material * dev_materials = NULL; static PathSegment * dev_paths = NULL; -static ShadeableIntersection * dev_intersections = NULL; +static ShadeableIntersection* dev_intersections = NULL; +static ShadeableIntersection* dev_cached_intersections = NULL; + // TODO: static variables for device memory, any extra info you need, etc // ... +struct RayIntersectionPredicate +{ + __host__ __device__ bool operator()(const ShadeableIntersection& shadeable_intersection) + { + + return (shadeable_intersection.t == -1); + } +}; + +struct RayTerminatePredicate +{ + RayTerminatePredicate() {} + + __host__ __device__ bool operator()(const PathSegment& path_segment) + { + + return (path_segment.remainingBounces > 0); + } +}; + +struct MaterialPredicate +{ + + MaterialPredicate() {} + + __host__ __device__ bool operator()(const ShadeableIntersection& first, const ShadeableIntersection& second) + { + return (first.materialId < second.materialId); + } +}; + void pathtraceInit(Scene *scene) { hst_scene = scene; const Camera &cam = hst_scene->state.camera; @@ -95,7 +149,8 @@ void pathtraceInit(Scene *scene) { cudaMalloc(&dev_intersections, pixelcount * sizeof(ShadeableIntersection)); cudaMemset(dev_intersections, 0, pixelcount * sizeof(ShadeableIntersection)); - // TODO: initialize any extra device memeory you need + cudaMalloc(&dev_cached_intersections, pixelcount * sizeof(ShadeableIntersection)); + cudaMemset(dev_cached_intersections, 0, pixelcount * sizeof(ShadeableIntersection)); checkCUDAError("pathtraceInit"); } @@ -106,6 +161,7 @@ void pathtraceFree() { cudaFree(dev_geoms); cudaFree(dev_materials); cudaFree(dev_intersections); + cudaFree(dev_cached_intersections); // TODO: clean up any extra device memory you created checkCUDAError("pathtraceFree"); @@ -121,24 +177,32 @@ void pathtraceFree() { */ __global__ void generateRayFromCamera(Camera cam, int iter, int traceDepth, PathSegment* pathSegments) { - int x = (blockIdx.x * blockDim.x) + threadIdx.x; - int y = (blockIdx.y * blockDim.y) + threadIdx.y; + const int x = (blockIdx.x * blockDim.x) + threadIdx.x; + const int y = (blockIdx.y * blockDim.y) + threadIdx.y; - if (x < cam.resolution.x && y < cam.resolution.y) { - int index = x + (y * cam.resolution.x); - PathSegment & segment = pathSegments[index]; + if (x < cam.resolution.x && y < cam.resolution.y) + { + const int index = x + (y * cam.resolution.x); + PathSegment& segment = pathSegments[index]; segment.ray.origin = cam.position; - segment.color = glm::vec3(1.0f, 1.0f, 1.0f); + segment.color = glm::vec3(1.0f, 1.0f, 1.0f); + + thrust::default_random_engine rng = makeSeededRandomEngine(iter, index, pathSegments[index].remainingBounces); + thrust::uniform_real_distribution u01(0, 2.f); + + const float jitterX = u01(rng);// (2.f * tanf(cam.fov.x / 2.f) * u1(rng) / cam.resolution.x); + const float jitterY = u01(rng);// (2.f * tanf(cam.fov.y / 2.f) * u2(rng) / cam.resolution.y); - // TODO: implement antialiasing by jittering the ray segment.ray.direction = glm::normalize(cam.view - - cam.right * cam.pixelLength.x * ((float)x - (float)cam.resolution.x * 0.5f) - - cam.up * cam.pixelLength.y * ((float)y - (float)cam.resolution.y * 0.5f) + - cam.right * ((cam.pixelLength.x * (jitterX + (float)x - (float)cam.resolution.x * 0.5f))) + - cam.up * cam.pixelLength.y * (jitterY + (float)y - (float)cam.resolution.y * 0.5f) ); segment.pixelIndex = index; segment.remainingBounces = traceDepth; + segment.isRayDead = false; + segment.isRefractedRay = false; } } @@ -152,62 +216,84 @@ __global__ void computeIntersections( , PathSegment * pathSegments , Geom * geoms , int geoms_size - , ShadeableIntersection * intersections + , ShadeableIntersection* intersections + , ShadeableIntersection* cacheIntersections ) { int path_index = blockIdx.x * blockDim.x + threadIdx.x; if (path_index < num_paths) { - PathSegment pathSegment = pathSegments[path_index]; + PathSegment& pathSegment = pathSegments[path_index]; float t; - glm::vec3 intersect_point; - glm::vec3 normal; float t_min = FLT_MAX; int hit_geom_index = -1; - bool outside = true; + const bool refractedRay = pathSegment.isRefractedRay; - glm::vec3 tmp_intersect; - glm::vec3 tmp_normal; + ShadeableIntersection intersection = intersections[path_index]; // naive parse through global geoms + // TODO : Clean this up + ShadeableIntersection tempIntersection = intersection; + ShadeableIntersection hitIntersection = tempIntersection; + for (int i = 0; i < geoms_size; i++) { - Geom & geom = geoms[i]; + Geom& geom = geoms[i]; + t = 0.0f; if (geom.type == CUBE) { - t = boxIntersectionTest(geom, pathSegment.ray, tmp_intersect, tmp_normal, outside); + t = Shapes::Cube::TestIntersection(&geom, pathSegment.ray, &tempIntersection, refractedRay); } else if (geom.type == SPHERE) { - t = sphereIntersectionTest(geom, pathSegment.ray, tmp_intersect, tmp_normal, outside); + t = Shapes::Sphere::TestIntersection(&geom, pathSegment.ray, &tempIntersection, refractedRay); + } + else if (geom.type == PLANE) + { + t = Shapes::SquarePlane::TestIntersection(&geom, pathSegment.ray, &tempIntersection, refractedRay); + } + else if (geom.type == IMPLICIT) + { + t = Shapes::Implicit::TestIntersection(&geom, &pathSegment.ray, &tempIntersection, refractedRay); } // TODO: add more intersection tests here... triangle? metaball? CSG? - // Compute the minimum t from the intersection tests to determine what - // scene geometry object was hit first. if (t > 0.0f && t_min > t) { t_min = t; hit_geom_index = i; - intersect_point = tmp_intersect; - normal = tmp_normal; + hitIntersection = tempIntersection; } } if (hit_geom_index == -1) { intersections[path_index].t = -1.0f; + intersections[path_index].m_didIntersect = false; } else { //The ray hits something intersections[path_index].t = t_min; intersections[path_index].materialId = geoms[hit_geom_index].materialid; - intersections[path_index].surfaceNormal = normal; + + intersections[path_index].m_surfaceNormal = hitIntersection.m_surfaceNormal; + intersections[path_index].m_surfaceTangent = hitIntersection.m_surfaceTangent; + intersections[path_index].m_surfaceBiTangent = hitIntersection.m_surfaceBiTangent; + + intersections[path_index].m_didIntersect = true; + intersections[path_index].m_intersectionPointWorld = hitIntersection.m_intersectionPointWorld; + intersections[path_index].m_tangentToWorld = hitIntersection.m_tangentToWorld; + intersections[path_index].m_worldToTangent = hitIntersection.m_worldToTangent; + } + + if(depth == 0) + { + cacheIntersections[path_index] = intersections[path_index]; } } } @@ -232,8 +318,8 @@ __global__ void shadeFakeMaterial ( int idx = blockIdx.x * blockDim.x + threadIdx.x; if (idx < num_paths) { - ShadeableIntersection intersection = shadeableIntersections[idx]; - if (intersection.t > 0.0f) { // if the intersection exists... + ShadeableIntersection& intersection = shadeableIntersections[idx]; + if (intersection.t > 0.0f && pathSegments[idx].remainingBounces > 0) { // if the intersection exists... // Set up the RNG // LOOK: this is how you use thrust's RNG! Please look at // makeSeededRandomEngine as well. @@ -251,7 +337,7 @@ __global__ void shadeFakeMaterial ( // like what you would expect from shading in a rasterizer like OpenGL. // TODO: replace this! you should be able to start with basically a one-liner else { - float lightTerm = glm::dot(intersection.surfaceNormal, glm::vec3(0.0f, 1.0f, 0.0f)); + float lightTerm = glm::dot(intersection.m_surfaceNormal, glm::vec3(0.0f, 1.0f, 0.0f)); pathSegments[idx].color *= (materialColor * lightTerm) * 0.3f + ((1.0f - intersection.t * 0.02f) * materialColor) * 0.7f; pathSegments[idx].color *= u01(rng); // apply some noise because why not } @@ -265,15 +351,195 @@ __global__ void shadeFakeMaterial ( } } +/* + * NaiveIntegratorShader + */ +__global__ void NaiveIntegratorShader( + int iter + , int num_paths + , ShadeableIntersection* shadeableIntersections + , PathSegment* pathSegments + , Material* materials +) +{ + int idx = blockIdx.x * blockDim.x + threadIdx.x; + + if (idx < num_paths) + { + ShadeableIntersection& intersection = shadeableIntersections[idx]; + + if (intersection.t > 0.0f) + { + + thrust::default_random_engine rng = makeSeededRandomEngine(iter, idx, pathSegments[idx].remainingBounces); + thrust::uniform_real_distribution u01(0, 1); + + Material material = materials[intersection.materialId]; + const glm::vec3 materialColor = material.color; + + glm::vec3 finalColor = (materialColor * material.emittance); + + // If the material indicates that the object was a light, "light" the ray + if(material.emittance > 0.0f || pathSegments[idx].remainingBounces <= 0) + { + finalColor *= pathSegments[idx].color; + + pathSegments[idx].color = finalColor; + pathSegments[idx].remainingBounces = 0; + pathSegments[idx].isRayDead = true; + } + else + { + // 1. Calculate WoW + const glm::vec3 woW = -pathSegments[idx].ray.direction; + glm::vec3 wiW(0.f); + glm::vec2 xi(u01(rng), u01(rng)); + + float pdf = 1.f; + + // 2. Get the wiw and pdf from the given material + const glm::vec3 sample_f_color = BSDF::Sample_F(woW, &wiW, &pdf, &xi, &material, &intersection); + + pathSegments[idx].isRefractedRay = material.hasRefractive; + + if(pdf < PDF_EPSILON) + { + pathSegments[idx].color = glm::vec3(1.f); + pathSegments[idx].remainingBounces = 0; + pathSegments[idx].isRayDead = true; + return; + } + + const float dotProduct = glm::dot(wiW, intersection.m_surfaceNormal); + const float lambertsTerm = glm::abs(dotProduct); + + // 3. Add the color to the path sement + // color *= (sample_f * lamberts) / pdf + pathSegments[idx].color *= (sample_f_color * lambertsTerm) / pdf; + + // 4. Update the ray direction and remove one bounce from path segment + const glm::vec3 originOffset = intersection.m_surfaceNormal * RAY_EPSILON * (dotProduct < 0 ? -1.f : 1.f); + + pathSegments[idx].ray.origin = intersection.m_intersectionPointWorld + originOffset; + pathSegments[idx].ray.direction = wiW; + pathSegments[idx].remainingBounces--; + } + } + else + { + pathSegments[idx].color = glm::vec3(0.0f); + pathSegments[idx].remainingBounces = 0; + pathSegments[idx].isRayDead = true; + } + } +} + +/* +* DirectLightingShader +*/ +__global__ void DirectLightingShader( + int iter + , int num_paths + , int num_geoms + , int num_lights + , ShadeableIntersection* shadeableIntersections + , PathSegment* pathSegments + , Geom* allGeometry + , Geom* lightGeometry + , Material* materials +) +{ + int idx = blockIdx.x * blockDim.x + threadIdx.x; + + if (idx < num_paths) + { + ShadeableIntersection& intersection = shadeableIntersections[idx]; + + if (intersection.t > 0.0f) + { + thrust::default_random_engine rng = makeSeededRandomEngine(iter, idx, pathSegments[idx].remainingBounces); + thrust::uniform_real_distribution u01(0, 1); + + Material material = materials[intersection.materialId]; + const glm::vec3 materialColor = material.color; + + glm::vec3 finalColor = (materialColor * material.emittance); + + // If the material indicates that the object was a light, "light" the ray + if(material.emittance > 0.0f || pathSegments[idx].remainingBounces <= 0 || num_lights == 0) + { + finalColor *= pathSegments[idx].color; + + pathSegments[idx].color = finalColor; + pathSegments[idx].remainingBounces = 0; + pathSegments[idx].isRayDead = true; + } + else + { + // 1. Initialize all the shit + const glm::vec3 woW = -pathSegments[idx].ray.direction; + glm::vec3 wiW_light(0.f); + glm::vec3 wiW_bsdf(0.f); + + glm::vec2 xi_light(u01(rng), u01(rng)); + glm::vec2 xi_bsdf(u01(rng), u01(rng)); + + int randomLight = int(u01(rng) * num_lights); + + float pdf_light = 1.f; + float pdf_bsdf = 1.f; + + Geom random_light = lightGeometry[randomLight]; + + + /*// 2. Get the wiw and pdf from the given material + const glm::vec3 sample_f_color = BSDF::Sample_F(woW, &wiW, &pdf, &xi, &material, &intersection); + + pathSegments[idx].isRefractedRay = material.hasRefractive; + + if(pdf < PDF_EPSILON) + { + pathSegments[idx].color = glm::vec3(1.f); + pathSegments[idx].remainingBounces = 0; + pathSegments[idx].isRayDead = true; + return; + } + + const float dotProduct = glm::dot(wiW, intersection.m_surfaceNormal); + const float lambertsTerm = glm::abs(dotProduct); + + // 3. Add the color to the path sement + // color *= (sample_f * lamberts) / pdf + pathSegments[idx].color *= (sample_f_color * lambertsTerm) / pdf; + + // 4. Update the ray direction and remove one bounce from path segment + const glm::vec3 originOffset = intersection.m_surfaceNormal * RAY_EPSILON * (dotProduct < 0 ? -1.f : 1.f); + + pathSegments[idx].ray.origin = intersection.m_intersectionPointWorld + originOffset; + pathSegments[idx].ray.direction = wiW; + pathSegments[idx].remainingBounces--;*/ + } + } + else + { + pathSegments[idx].color = glm::vec3(0.0f); + pathSegments[idx].remainingBounces = 0; + pathSegments[idx].isRayDead = true; + } + } +} + + + // Add the current iteration's output to the overall image -__global__ void finalGather(int nPaths, glm::vec3 * image, PathSegment * iterationPaths) +__global__ void finalGather(int nPaths, int totalIterations, glm::vec3* image, PathSegment* iterationPaths) { int index = (blockIdx.x * blockDim.x) + threadIdx.x; if (index < nPaths) { PathSegment iterationPath = iterationPaths[index]; - image[iterationPath.pixelIndex] += iterationPath.color; + image[iterationPath.pixelIndex] += iterationPath.color;// / float(totalIterations)); } } @@ -281,7 +547,7 @@ __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(uchar4 *pbo, int frame, int iter) { +void pathtrace(uchar4 *pbo, int frame, int iter, int totalIterations) { const int traceDepth = hst_scene->state.traceDepth; const Camera &cam = hst_scene->state.camera; const int pixelcount = cam.resolution.x * cam.resolution.y; @@ -295,90 +561,71 @@ void pathtrace(uchar4 *pbo, int frame, int iter) { // 1D block for path tracing const int blockSize1d = 128; - /////////////////////////////////////////////////////////////////////////// - - // Recap: - // * Initialize array of path rays (using rays that come out of the camera) - // * You can pass the Camera object to that kernel. - // * Each path ray must carry at minimum a (ray, color) pair, - // * where color starts as the multiplicative identity, white = (1, 1, 1). - // * This has already been done for you. - // * For each depth: - // * Compute an intersection in the scene for each path ray. - // A very naive version of this has been implemented for you, but feel - // free to add more primitives and/or a better algorithm. - // Currently, intersection distance is recorded as a parametric distance, - // t, or a "distance along the ray." t = -1.0 indicates no intersection. - // * Color is attenuated (multiplied) by reflections off of any object - // * TODO: Stream compact away all of the terminated paths. - // You may use either your implementation or `thrust::remove_if` or its - // cousins. - // * Note that you can't really use a 2D kernel launch any more - switch - // to 1D. - // * TODO: Shade the rays that intersected something or didn't bottom out. - // That is, color the ray by performing a color computation according - // to the shader, then generate a new ray to continue the ray path. - // We recommend just updating the ray's PathSegment in place. - // Note that this step may come before or after stream compaction, - // since some shaders you write may also cause a path to terminate. - // * Finally, add this iteration's results to the image. This has been done - // for you. - - // TODO: perform one iteration of path tracing - - generateRayFromCamera <<>>(cam, iter, traceDepth, dev_paths); + generateRayFromCamera <<>>(cam, iter, traceDepth, dev_paths); checkCUDAError("generate camera ray"); int depth = 0; PathSegment* dev_path_end = dev_paths + pixelcount; - int num_paths = dev_path_end - dev_paths; + const int num_paths = dev_path_end - dev_paths; + int curr_paths = num_paths; // --- PathSegment Tracing Stage --- // Shoot ray into scene, bounce between objects, push shading chunks - bool iterationComplete = false; - while (!iterationComplete) { - - // clean shading chunks - cudaMemset(dev_intersections, 0, pixelcount * sizeof(ShadeableIntersection)); - - // tracing - dim3 numblocksPathSegmentTracing = (num_paths + blockSize1d - 1) / blockSize1d; - computeIntersections <<>> ( - depth - , num_paths - , dev_paths - , dev_geoms - , hst_scene->geoms.size() - , dev_intersections - ); - checkCUDAError("trace one bounce"); - cudaDeviceSynchronize(); - depth++; - - - // TODO: - // --- Shading Stage --- - // Shade path segments based on intersections and generate new rays by - // evaluating the BSDF. - // Start off with just a big kernel that handles all the different - // materials you have in the scenefile. - // TODO: compare between directly shading the path segments and shading - // path segments that have been reshuffled to be contiguous in memory. - - shadeFakeMaterial<<>> ( - iter, - num_paths, - dev_intersections, - dev_paths, - dev_materials - ); - iterationComplete = true; // TODO: should be based off stream compaction results. + bool iterationComplete = false; + while (!iterationComplete) + { + // clean shading chunks + cudaMemset(dev_intersections, 0, pixelcount * sizeof(ShadeableIntersection)); + + // tracing + dim3 numblocksPathSegmentTracing = (curr_paths + blockSize1d - 1) / blockSize1d; + + const bool useCachedIntersection = (depth == 0 && iter != 1); + + if(!useCachedIntersection) + { + computeIntersections <<>> ( + depth + , curr_paths + , dev_paths + , dev_geoms + , hst_scene->geoms.size() + , dev_intersections + , dev_cached_intersections + ); + checkCUDAError("trace one bounce"); + cudaDeviceSynchronize(); + } + depth++; + + // Sort by Material + thrust::sort_by_key(thrust::device, dev_intersections, dev_intersections + curr_paths, dev_paths, MaterialPredicate()); + + // 1. Do stream and remove rays that have no intersection. + // Not needed for now + + // 2. Perform Color calculation using BSDF for Naive. + NaiveIntegratorShader << < numblocksPathSegmentTracing, blockSize1d >> > ( + iter, + curr_paths, + (useCachedIntersection ? dev_cached_intersections : dev_intersections), + dev_paths, + dev_materials + ); + + // 3. Remove any rays that have reached maximum bounces. + dev_path_end = thrust::partition(thrust::device, dev_paths, dev_paths + curr_paths, RayTerminatePredicate()); + cudaDeviceSynchronize(); + curr_paths = (dev_path_end - dev_paths); + + // This should be based on result of (3). + iterationComplete = (depth > traceDepth || curr_paths <= 0); } - // Assemble this iteration and apply it to the image - dim3 numBlocksPixels = (pixelcount + blockSize1d - 1) / blockSize1d; - finalGather<<>>(num_paths, dev_image, dev_paths); + // Assemble this iteration and apply it to the image + dim3 numBlocksPixels = (pixelcount + blockSize1d - 1) / blockSize1d; + finalGather<<>>(pixelcount, totalIterations, dev_image, dev_paths); /////////////////////////////////////////////////////////////////////////// diff --git a/src/pathtrace.h b/src/pathtrace.h index 1241227..b2c16f4 100644 --- a/src/pathtrace.h +++ b/src/pathtrace.h @@ -5,4 +5,4 @@ void pathtraceInit(Scene *scene); void pathtraceFree(); -void pathtrace(uchar4 *pbo, int frame, int iteration); +void pathtrace(uchar4 *pbo, int frame, int iteration, int totalIterations); diff --git a/src/scene.cpp b/src/scene.cpp index cbae043..846e67e 100644 --- a/src/scene.cpp +++ b/src/scene.cpp @@ -51,7 +51,13 @@ int Scene::loadGeom(string objectid) { } else if (strcmp(line.c_str(), "cube") == 0) { cout << "Creating new cube..." << endl; newGeom.type = CUBE; - } + } else if (strcmp(line.c_str(), "plane") == 0) { + cout << "Creating new plane..." << endl; + newGeom.type = PLANE; + } else if (strcmp(line.c_str(), "implicit") == 0) { + cout << "Creating new implicit..." << endl; + newGeom.type = IMPLICIT; + } } //link material @@ -150,39 +156,85 @@ int Scene::loadCamera() { return 1; } -int Scene::loadMaterial(string materialid) { +int Scene::loadMaterial(string materialid) +{ int id = atoi(materialid.c_str()); - if (id != materials.size()) { + if (id != materials.size()) + { cout << "ERROR: MATERIAL ID does not match expected number of materials" << endl; return -1; - } else { + } + else + { cout << "Loading Material " << id << "..." << endl; Material newMaterial; + newMaterial.numBxdfs = 0; //load static properties - for (int i = 0; i < 7; i++) { + for (int i = 0; i < 8; i++) + { string line; utilityCore::safeGetline(fp_in, line); vector tokens = utilityCore::tokenizeString(line); - if (strcmp(tokens[0].c_str(), "RGB") == 0) { + + if (strcmp(tokens[0].c_str(), "RGB") == 0) + { glm::vec3 color( atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str()) ); newMaterial.color = color; - } else if (strcmp(tokens[0].c_str(), "SPECEX") == 0) { + } + else if (strcmp(tokens[0].c_str(), "SPECEX") == 0) + { newMaterial.specular.exponent = atof(tokens[1].c_str()); - } else if (strcmp(tokens[0].c_str(), "SPECRGB") == 0) { + const bool isSpecular = newMaterial.specular.exponent != 0; + if(isSpecular) + { + newMaterial.bxdfTypes[newMaterial.numBxdfs++] = BxDFType::BSDF_SPECULAR; + } + } + else if (strcmp(tokens[0].c_str(), "SPECRGB") == 0) + { glm::vec3 specColor(atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str())); newMaterial.specular.color = specColor; - } else if (strcmp(tokens[0].c_str(), "REFL") == 0) { + } + else if (strcmp(tokens[0].c_str(), "REFL") == 0) + { newMaterial.hasReflective = atof(tokens[1].c_str()); - } else if (strcmp(tokens[0].c_str(), "REFR") == 0) { + const bool isReflective = newMaterial.hasReflective != 0; + if(isReflective) + { + newMaterial.bxdfTypes[newMaterial.numBxdfs++] = BxDFType::BSDF_REFLECTION; + } + } + else if (strcmp(tokens[0].c_str(), "REFR") == 0) + { newMaterial.hasRefractive = atof(tokens[1].c_str()); - } else if (strcmp(tokens[0].c_str(), "REFRIOR") == 0) { + newMaterial.refractive.etaA = 1.f; + newMaterial.refractive.etaB = 1.52f; + + const bool isRefractive = newMaterial.hasRefractive != 0; + if(isRefractive) + { + newMaterial.bxdfTypes[newMaterial.numBxdfs++] = BxDFType::BSDF_TRANSMISSION; + } + } + else if (strcmp(tokens[0].c_str(), "REFRIOR") == 0) + { newMaterial.indexOfRefraction = atof(tokens[1].c_str()); - } else if (strcmp(tokens[0].c_str(), "EMITTANCE") == 0) { + } + else if (strcmp(tokens[0].c_str(), "EMITTANCE") == 0) + { newMaterial.emittance = atof(tokens[1].c_str()); } + else if (strcmp(tokens[0].c_str(), "DIFFUSE") == 0) + { + const bool isDiffuse = atof(tokens[1].c_str()) != 0; + newMaterial.bxdfTypes[newMaterial.numBxdfs++] = BxDFType::BSDF_DIFFUSE; + } + } + materials.push_back(newMaterial); return 1; } } + diff --git a/src/sceneStructs.cu b/src/sceneStructs.cu new file mode 100644 index 0000000..bee9200 --- /dev/null +++ b/src/sceneStructs.cu @@ -0,0 +1,184 @@ +#pragma once + +#include +#include +#include + +#include "glm/glm.hpp" + +#include "common.cu" +#include "warpfunctions.cu" + +#define BACKGROUND_COLOR (glm::vec3(0.0f)) + +#define MAX_BXDFS 3 + +enum BxDFType +{ + BSDF_REFLECTION = 1 << 0, // This BxDF handles rays that are reflected off surfaces + BSDF_TRANSMISSION = 1 << 1, // This BxDF handles rays that are transmitted through surfaces + BSDF_DIFFUSE = 1 << 2, // This BxDF represents diffuse energy scattering, which is uniformly random + BSDF_GLOSSY = 1 << 3, // This BxDF represents glossy energy scattering, which is biased toward certain directions + BSDF_SPECULAR = 1 << 4, // This BxDF handles specular energy scattering, which has no element of randomness + BSDF_ALL = BSDF_DIFFUSE | BSDF_GLOSSY | BSDF_SPECULAR | BSDF_REFLECTION | BSDF_TRANSMISSION +}; + + +enum GeomType +{ + SPHERE, + CUBE, +}; + +struct Ray { + glm::vec3 origin; + glm::vec3 direction; +}; + +struct Geom +{ + enum GeomType type; + int materialid; + glm::vec3 translation; + glm::vec3 rotation; + glm::vec3 scale; + glm::mat4 transform; + glm::mat4 inverseTransform; + glm::mat4 invTranspose; +}; + +//--------------------------------------------------------------------------------- +// BxDF +//--------------------------------------------------------------------------------- +struct Bxdf +{ + __device__ virtual glm::vec3 SampleF(const glm::vec3* wo, glm::vec3* wi, float* pdf, const glm::vec2* xi, const glm::vec3 color); + __device__ virtual float Pdf(const glm::vec3* wo, glm::vec3* wi); + __device__ virtual glm::vec3 F(const glm::vec3* wo, const glm::vec3* wi, const glm::vec3 color); +}; + +//--------------------------------------------------------------------------------- +// Lamberts or Diffuse BRDF +//--------------------------------------------------------------------------------- +struct LambertsBrdf : public Bxdf +{ + // Lamberts BRDF ------------------------------------------------------- + __device__ glm::vec3 F(const glm::vec3* wo, const glm::vec3* wi, const glm::vec3 color) + { + return color * InvPi; + } + + __device__ float Pdf(const glm::vec3* wo, glm::vec3* wi) + { + return Common::SameHemisphere(wo, wi) ? Common::AbsCosTheta(wi) * InvPi : 0; + } + + __device__ glm::vec3 SampleF(const glm::vec3* wo, glm::vec3* wi, float* pdf, const glm::vec2* xi, const glm::vec3 color) + { + // 1. Cosine sample the hemisphere + *wi = WarpFunctions::SquareToHemisphereCosine(xi); + + if (wo->z < 0) { + wi->z *= -1; + } + + // 2. Calculate the pdf + *pdf = Pdf(wo, wi); + + // 3. return f + return F(wo, wi, color); + } +}; + +//--------------------------------------------------------------------------------- +// Specular BRDF +//--------------------------------------------------------------------------------- +struct SpecularBrdf : public Bxdf +{ + // Specular BRDF ------------------------------------------------------- + __device__ glm::vec3 F(const glm::vec3* wo, const glm::vec3* wi, const glm::vec3 color) + { + return color; + } + + __device__ float Pdf(const glm::vec3* wo, glm::vec3* wi) + { + return 1.f; + } + + __device__ glm::vec3 SampleF(const glm::vec3* wo, glm::vec3* wi, float* pdf, const glm::vec2* xi, const glm::vec3 color) + { + *wi = glm::vec3(-(*wo).x, -(*wo).y, (*wo).z); + *pdf = Pdf(wo, wi); + return F(wo, wi, color); + } +}; + +struct Material +{ + glm::vec3 color; + struct { + float exponent; + glm::vec3 color; + } specular; + float hasReflective; + float hasRefractive; + float indexOfRefraction; + float emittance; + + Bxdf bxdfs[MAX_BXDFS]; + int numBxdfs = 0; + + void AddBxdf(BxDFType bxdfType) + { + if(bxdfType == BSDF_DIFFUSE && numBxdfs < (MAX_BXDFS - 1)) + { + bxdfs[numBxdfs++] = LambertsBrdf(); + } + } +}; + +struct Camera { + glm::ivec2 resolution; + glm::vec3 position; + glm::vec3 lookAt; + glm::vec3 view; + glm::vec3 up; + glm::vec3 right; + glm::vec2 fov; + glm::vec2 pixelLength; +}; + +struct RenderState { + Camera camera; + unsigned int iterations; + int traceDepth; + std::vector image; + std::string imageName; +}; + +struct PathSegment { + Ray ray; + glm::vec3 color; + int pixelIndex; + int remainingBounces; + bool isRayDead; +}; + +// Use with a corresponding PathSegment to do: +// 1) color contribution computation +// 2) BSDF evaluation: generate a new ray +struct ShadeableIntersection +{ + glm::vec3 m_intersectionPointWorld; + bool m_didIntersect; + float t; + glm::vec3 m_surfaceNormal; + glm::vec3 m_surfaceTangent; + glm::vec3 m_surfaceBiTangent; + + glm::mat3 m_worldToTangent; + glm::mat3 m_tangentToWorld; + + int materialId; +}; diff --git a/src/sceneStructs.h b/src/sceneStructs.h index b38b820..d25e089 100644 --- a/src/sceneStructs.h +++ b/src/sceneStructs.h @@ -7,9 +7,24 @@ #define BACKGROUND_COLOR (glm::vec3(0.0f)) -enum GeomType { +#define MAX_BXDFS 3 + +enum BxDFType +{ + BSDF_REFLECTION = 1 << 0, // This BxDF handles rays that are reflected off surfaces + BSDF_TRANSMISSION = 1 << 1, // This BxDF handles rays that are transmitted through surfaces + BSDF_DIFFUSE = 1 << 2, // This BxDF represents diffuse energy scattering, which is uniformly random + BSDF_GLOSSY = 1 << 3, // This BxDF represents glossy energy scattering, which is biased toward certain directions + BSDF_SPECULAR = 1 << 4, // This BxDF handles specular energy scattering, which has no element of randomness + BSDF_ALL = BSDF_DIFFUSE | BSDF_GLOSSY | BSDF_SPECULAR | BSDF_REFLECTION | BSDF_TRANSMISSION +}; + +enum GeomType +{ SPHERE, CUBE, + PLANE, + IMPLICIT }; struct Ray { @@ -17,7 +32,8 @@ struct Ray { glm::vec3 direction; }; -struct Geom { +struct Geom +{ enum GeomType type; int materialid; glm::vec3 translation; @@ -28,16 +44,27 @@ struct Geom { glm::mat4 invTranspose; }; -struct Material { + +struct Material +{ glm::vec3 color; struct { float exponent; glm::vec3 color; } specular; + + struct { + float etaA; + float etaB; + } refractive; + float hasReflective; float hasRefractive; float indexOfRefraction; float emittance; + + BxDFType bxdfTypes[MAX_BXDFS]; + int numBxdfs; }; struct Camera { @@ -59,18 +86,31 @@ struct RenderState { std::string imageName; }; -struct PathSegment { +struct PathSegment +{ Ray ray; glm::vec3 color; int pixelIndex; int remainingBounces; + bool isRayDead; + bool isRefractedRay; }; // Use with a corresponding PathSegment to do: // 1) color contribution computation // 2) BSDF evaluation: generate a new ray -struct ShadeableIntersection { - float t; - glm::vec3 surfaceNormal; - int materialId; +struct ShadeableIntersection +{ + glm::vec3 m_intersectionPointWorld; + bool m_didIntersect; + float t; + glm::vec3 m_surfaceNormal; + glm::vec3 m_surfaceTangent; + glm::vec3 m_surfaceBiTangent; + + glm::mat3 m_worldToTangent; + glm::mat3 m_tangentToWorld; + + int materialId; }; + diff --git a/src/shapes.cu b/src/shapes.cu new file mode 100644 index 0000000..0c6bf18 --- /dev/null +++ b/src/shapes.cu @@ -0,0 +1,415 @@ +#pragma once + +#include +#include "glm/glm.hpp" + +namespace Shapes +{ + /** + * Multiplies a mat4 and a vec4 and returns a vec3 clipped from the vec4. + */ + __host__ __device__ glm::vec3 multiplyMV(glm::mat4 m, glm::vec4 v) + { + return glm::vec3(m * v); + } + + namespace SquarePlane + { + __host__ __device__ float Area(const Geom* plane) + { + return (plane->scale.x * plane->scale.y); + } + + __host__ __device__ float TestIntersection(const Geom* plane, Ray r, ShadeableIntersection* intersection, bool outside) + { + const glm::vec3 ray_origin = multiplyMV(plane->inverseTransform, glm::vec4(r.origin, 1.0f)); + const glm::vec3 ray_direction = glm::normalize(multiplyMV(plane->inverseTransform, glm::vec4(r.direction, 0.0f))); + + const float t = (glm::vec3(0.5f, 0.5f, 0) - ray_origin).z / ray_direction.z; + const glm::vec3 P = glm::vec3(t * ray_direction + ray_origin); + + //Check that P is within the bounds of the square + if(t > 0 && P.x >= -0.5f && P.x <= 0.5f && P.y >= -0.5f && P.y <= 0.5f) + { + const glm::vec3 objSpaceIntersection = P; + const glm::vec3 intersectionPoint = multiplyMV(plane->transform, glm::vec4(objSpaceIntersection, 1.0f)); + intersection->m_intersectionPointWorld = intersectionPoint; + + // Computing normal, tangent and bitangent + const glm::vec3 normal = glm::normalize(glm::mat3(plane->inverseTransform) * glm::vec3(0.f, 0.f, 1.f)); + const glm::vec3 tangent = glm::normalize(glm::vec3(plane->transform * glm::vec4(1.f, 0.f, 0.f, 0.f))); + const glm::vec3 bitangent = glm::normalize(glm::vec3(plane->transform * glm::vec4(0.f, 1.f, 0.f, 0.f))); + + intersection->m_surfaceNormal = normal; + intersection->m_surfaceTangent = tangent; + intersection->m_surfaceBiTangent = bitangent; + + // Compute transformation matrices + intersection->m_tangentToWorld = glm::mat3(tangent, bitangent, normal); + intersection->m_worldToTangent = glm::transpose(intersection->m_tangentToWorld); + + return glm::length(r.origin - intersectionPoint); + + } + return -1; + } + + __host__ __device__ void Sample(const Geom* plane, const glm::vec2* xi, float *pdf, glm::vec3* intersectionPoint, glm::vec3* intersectionNormal) + { + *intersectionPoint = glm::vec3(plane->transform * glm::vec4((*xi)[0] - 0.5f, (*xi)[1] - 0.5f, 0.f, 1.f)); + *intersectionNormal = glm::normalize(glm::mat3(plane->inverseTransform) * glm::vec3(0.f, 0.f, 1.f)); + *pdf = 1.f / Area(plane); + } + } + + namespace Cube + { + namespace + { + __host__ __device__ int GetFaceIndex(const glm::vec3& P) + { + int idx = 0; + float val = -1; + for(int i = 0; i < 3; i++){ + if(glm::abs(P[i]) > val){ + idx = i * glm::sign(P[i]); + val = glm::abs(P[i]); + } + } + return idx; + } + + __host__ __device__ glm::vec3 GetCubeNormal(const glm::vec3& P, const int faceIndex) + { + glm::vec3 N(0,0,0); + N[faceIndex] = glm::sign(P[faceIndex]); + return N; + } + + __host__ __device__ glm::vec3 GetCubeTangent(const glm::vec3& P, const int faceIndex) + { + glm::vec3 N(0,0,0); + if(faceIndex == 0 && glm::sign(P[faceIndex]) > 0) { + N = glm::vec3(0.f, 0.f, -1.f); + } else if(faceIndex == 1 && glm::sign(P[faceIndex]) > 0) { + N = glm::vec3(0.f, 0.f, 1.f); + } else if(faceIndex == 2 && glm::sign(P[faceIndex]) > 0) { + N = glm::vec3(1.f, 0.f, 0.f); + } else if(faceIndex == 0 && glm::sign(P[faceIndex]) < 0) { + N = glm::vec3(0.f, 0.f, 1.f); + } else if(faceIndex == 1 && glm::sign(P[faceIndex]) < 0) { + N = glm::vec3(1.f, 0.f, 0.f); + } else if(faceIndex == 2 && glm::sign(P[faceIndex]) < 0) { + N = glm::vec3(-1.f, 0.f, 0.f); + } + return N; + } + + } // Anonymous namespace end + + /** + * Test intersection between a ray and a transformed cube. Untransformed, + * the cube ranges from -0.5 to 0.5 in each axis and is centered at the origin. + * + * @param intersectionPoint Output parameter for point of intersection. + * @param normal Output parameter for surface normal. + * @param outside Output param for whether the ray came from outside. + * @return Ray parameter `t` value. -1 if no intersection. + */ + __host__ __device__ float TestIntersection(const Geom* box, Ray r, ShadeableIntersection* intersection, bool outside) + { + glm::vec3 ray_origin = multiplyMV(box->inverseTransform, glm::vec4(r.origin , 1.0f)); + glm::vec3 ray_direction = glm::normalize(multiplyMV(box->inverseTransform, glm::vec4(r.direction, 0.0f))); + + float t_n = -1000000; + float t_f = 1000000; + + for(int i = 0; i < 3; i++) + { + //Ray parallel to slab check + if(ray_direction[i] == 0) + { + if(ray_origin[i] < -0.5f || ray_origin[i] > 0.5f) + { + return -1; + } + } + //If not parallel, do slab intersect check + float t0 = (-0.5f - ray_origin[i])/ray_direction[i]; + float t1 = (0.5f - ray_origin[i])/ray_direction[i]; + if(t0 > t1) + { + float temp = t1; + t1 = t0; + t0 = temp; + } + if(t0 > t_n) + { + t_n = t0; + } + + if(t1 < t_f) + { + t_f = t1; + } + + } + if(t_n < t_f) + { + const float t = t_n > 0 ? t_n : t_f; + + if (t < 0) + { + return -1; + } + //Lastly, transform the point found in object space by T + const glm::vec3 objSpaceIntersection = ray_origin + t * ray_direction; + const glm::vec3 intersectionPoint = multiplyMV(box->transform, glm::vec4(objSpaceIntersection, 1.0f)); + intersection->m_intersectionPointWorld = intersectionPoint; + + // Computing normal, tangent and bitangent + const int faceIndex = glm::abs(GetFaceIndex(objSpaceIntersection)); + const glm::vec3 localNormal = GetCubeNormal(objSpaceIntersection, faceIndex); + const glm::vec3 localTangent = GetCubeTangent(objSpaceIntersection, faceIndex); + + const glm::vec3 normal = glm::normalize(multiplyMV(box->transform, glm::vec4(localNormal, 0.0f))); + const glm::vec3 tangent = glm::normalize(glm::vec3(box->transform * glm::vec4(localTangent, 0.0f))); + const glm::vec3 bitangent = glm::normalize(glm::cross(normal, tangent)); + + intersection->m_surfaceNormal = normal; + intersection->m_surfaceTangent = tangent; + intersection->m_surfaceBiTangent = bitangent; + + // Compute transformation matrices + intersection->m_tangentToWorld = glm::mat3(tangent, bitangent, normal); + intersection->m_worldToTangent = glm::transpose(intersection->m_tangentToWorld); + + return glm::length(r.origin - intersectionPoint); + } + return -1; + } + } + + namespace Sphere + { + /** + * Test intersection between a ray and a transformed sphere. Untransformed, + * the sphere always has radius of 0.5 and is centered at the origin. + * + * @param intersectionPoint Output parameter for point of intersection. + * @param normal Output parameter for surface normal. + * @param outside Output param for whether the ray came from outside. + * @return Ray parameter `t` value. -1 if no intersection. + */ + __host__ __device__ float TestIntersection(Geom* sphere, Ray r, ShadeableIntersection* intersection, bool outside) + { + glm::vec3 ro = multiplyMV(sphere->inverseTransform, glm::vec4(r.origin, 1.0f)); + glm::vec3 rd = glm::normalize(multiplyMV(sphere->inverseTransform, glm::vec4(r.direction, 0.0f))); + + float A = pow(rd.x, 2.f) + pow(rd.y, 2.f) + pow(rd.z, 2.f); + float B = 2*(rd.x*ro.x + rd.y * ro.y + rd.z * ro.z); + float C = pow(ro.x, 2.f) + pow(ro.y, 2.f) + pow(ro.z, 2.f) - 0.25f;//Radius is 0.5f + float discriminant = B*B - 4*A*C; + + //If the discriminant is negative, then there is no real root + if(discriminant < 0) + { + return -1; + } + + float sqrtDiscriminant = sqrt(discriminant); + + float t = (-B - sqrtDiscriminant)/(2*A); + + if(t < 0) + { + t = (-B + sqrtDiscriminant)/(2*A); + } + + if(t >= 0) + { + const glm::vec3 objspaceIntersection = ro + t*rd; + const glm::vec3 intersectionPoint = multiplyMV(sphere->transform, glm::vec4(objspaceIntersection, 1.f)); + + intersection->m_intersectionPointWorld = intersectionPoint; + + const glm::vec3 normal = glm::normalize(multiplyMV(sphere->invTranspose, glm::vec4(objspaceIntersection, 0.f))) * (outside ? -1.f : 1.f); + const glm::vec3 tangent = glm::normalize(glm::mat3(sphere->transform) * glm::cross(glm::vec3(0.f, 1.f, 0.f), glm::normalize(objspaceIntersection))); + const glm::vec3 bitangent = glm::normalize(glm::cross(normal, tangent)); + + intersection->m_surfaceNormal = normal; + intersection->m_surfaceTangent = tangent; + intersection->m_surfaceBiTangent = bitangent; + + // Compute transformation matrices + intersection->m_tangentToWorld = glm::mat3(tangent, bitangent, normal); + intersection->m_worldToTangent = glm::transpose(glm::mat3(tangent, bitangent, normal)); + + return glm::length(r.origin - intersectionPoint); + } + return -1; + } + } + + namespace Mesh + { + + } + + namespace Implicit + { + + #define SPHERE_RADIUS 1.f + #define RAY_MARCH_EPSILON 0.000001f + #define RAY_STEPS 100 + #define ROUNDED_BOX_RADIUS 0.1f + #define CYLINDER_RADIUS 0.5f + #define CYLINDER_HEIGHT 1.0f + #define BOX_WIDTH 1.0f + + __host__ __device__ float SphereSDF(glm::vec3 point) + { + return glm::length(point) - SPHERE_RADIUS; + } + + __host__ __device__ float BoxSDF(glm::vec3 point) + { + glm::vec3 d = glm::abs(point) - glm::vec3(BOX_WIDTH); + return glm::min(glm::max(d.x, glm::max(d.y,d.z)), 0.f) + glm::length(glm::max(d, glm::vec3(0.f))); + } + + + __host__ __device__ float UdRoundBox(glm::vec3 p) + { + return glm::length(glm::max(glm::abs(p)- glm::vec3(1.f), glm::vec3(0.0))) - ROUNDED_BOX_RADIUS; + } + + __host__ __device__ float TorusSDF(glm::vec3 p) + { + glm::vec2 t = glm::vec2(2.f, 1); + glm::vec2 q = glm::vec2(glm::length(glm::vec2(p.x, p.z)) - t.x, p.y); + return glm::length(q) - t.y; + } + + __host__ __device__ float SdCappedCylinder(glm::vec3 p) + { + glm::vec2 d = glm::abs(glm::vec2(glm::length(glm::vec2(p.x, p.z)), p.y)) - glm::vec2(CYLINDER_RADIUS, CYLINDER_HEIGHT); + return glm::min(glm::max(d.x,d.y), 0.f) + glm::length(glm::max(d, 0.f)); + } + + __host__ __device__ float FuncUnion(float d1, float d2) + { + return glm::max(d1, d2); + } + + __host__ __device__ float FuncIntersection(float d1, float d2) + { + return glm::min(d1, d2); + } + + __host__ __device__ float FuncSubtract(float d1, float d2) + { + return glm::max(-d1, d2); + } + + __host__ __device__ float FunkySDF(glm::vec3 p) + { + return 0.f; + } + + __host__ __device__ float Raymarch(glm::vec3* eye, glm::vec3* dir, float max) + { + float depth = 0.f; + for(int i = 0; i < RAY_STEPS; ++i) + { + const glm::vec3 point = *eye + depth * (*dir); + float dist = TorusSDF(point); + if(dist < RAY_MARCH_EPSILON) + { + return depth; + } + if(depth >= max) + { + return max; + } + depth += dist; + } + return max; + } + + __host__ __device__ void ComputeTBN(const glm::vec3 P, const Geom* geometry, glm::vec3* nor, glm::vec3* tan, glm::vec3* bit) + { + + const float epsilon = 0.0001f; + + float dx = TorusSDF(glm::vec3(P[0] + epsilon, P[1], P[2])) - TorusSDF(glm::vec3(P[0] - epsilon, P[1], P[2])); + float dy = TorusSDF(glm::vec3(P[0], P[1] + epsilon, P[2])) - TorusSDF(glm::vec3(P[0], P[1] - epsilon, P[2])); + float dz = TorusSDF(glm::vec3(P[0], P[1], P[2] + epsilon)) - TorusSDF(glm::vec3(P[0], P[1], P[2] - epsilon)); + + glm::vec3 normal = glm::normalize(glm::vec3(dx, dy, dz)); + + float tanX = P.x + epsilon; + float tanY = P.y + epsilon; + float tanZ = ((dx * (tanX - P.x) + dy * (tanY - P.y)) / (-dz)) + P.z; + + glm::vec3 tangent = glm::normalize(glm::vec3(tanX, tanY, tanZ)); + glm::vec3 bitangent = glm::normalize(glm::cross(normal, tangent)); + + *nor = glm::normalize(glm::mat3(geometry->invTranspose) * normal); + *tan = glm::normalize(glm::mat3(geometry->transform) * tangent); + *bit = glm::normalize(glm::mat3(geometry->transform) * bitangent); + } + + /** + * Test intersection between a ray and a transformed sphere. Untransformed, + * the sphere always has radius of 0.5 and is centered at the origin. + * + * @param intersectionPoint Output parameter for point of intersection. + * @param normal Output parameter for surface normal. + * @param outside Output param for whether the ray came from outside. + * @return Ray parameter `t` value. -1 if no intersection. + */ + __host__ __device__ float TestIntersection(Geom* geometry, const Ray* r, ShadeableIntersection* intersection, bool outside) + { + glm::vec3 ro = multiplyMV(geometry->inverseTransform, glm::vec4(r->origin, 1.0f)); + glm::vec3 rd = glm::normalize(multiplyMV(geometry->inverseTransform, glm::vec4(r->direction, 0.0f))); + + float max = 2000; + + float t = Raymarch(&ro, &rd, max); + if(t < max) + { + const glm::vec3 objspaceIntersection = ro + t*rd; + const glm::vec3 intersectionPoint = multiplyMV(geometry->transform, glm::vec4(objspaceIntersection, 1.f)); + + intersection->m_intersectionPointWorld = intersectionPoint; + + glm::vec3 normal(0.f); + glm::vec3 tangent(0.f); + glm::vec3 bitangent(0.f); + + ComputeTBN(objspaceIntersection, geometry, &normal, &tangent, &bitangent); + + intersection->m_surfaceNormal = normal; + intersection->m_surfaceTangent = tangent; + intersection->m_surfaceBiTangent = bitangent; + + // Compute transformation matrices + intersection->m_tangentToWorld = glm::mat3(tangent, bitangent, normal); + intersection->m_worldToTangent = glm::transpose(glm::mat3(tangent, bitangent, normal)); + + return glm::length(r->origin - intersectionPoint); + } + return -1; + } + } + + +} // namespace Shapes end + + + + + + + diff --git a/src/warpfunctions.cu b/src/warpfunctions.cu new file mode 100644 index 0000000..770aca7 --- /dev/null +++ b/src/warpfunctions.cu @@ -0,0 +1,129 @@ +#include +#include "sceneStructs.h" +#include "common.cu" + +namespace WarpFunctions +{ + + __host__ __device__ glm::vec3 SquareToDiskUniform(const glm::vec2* sample) + { + const float radius = (*sample)[0]; + glm::vec3 result; + const float angle = (*sample)[1] * 2.f * Pi; + result[0] = radius * cos(angle); + result[1] = radius * sin(angle); + result[2] = 0; + return result; + } + + __host__ __device__ glm::vec3 SquareToDiskConcentric(const glm::vec2* sample) + { + glm::vec3 result; + float phi, r; + float a = 2.f * (*sample)[0] - 1; + float b = 1 - 2.f * (*sample)[1]; + + if (a > -b) { + if (a > b) { + r = a; + phi = (PiOver4) * (b / a); + } + else { + r = b; + phi = (PiOver4) * (2.f - (a / b)); + } + } + else { + if (a < b) { + r = -a; + phi = (PiOver4) * (4.f + (b / a)); + } + else { + r = -b; + if (b != 0) { + phi = (PiOver4) * (6.f - (a / b)); + } + else { + phi = 0; + } + } + } + result[0] = r * cos(phi); + result[1] = r * sin(phi); + result[2] = 0; + return result; + } + + __host__ __device__ float SquareToDiskPDF(const glm::vec3* sample) + { + return InvPi; + } + + __host__ __device__ glm::vec3 SquareToSphereUniform(const glm::vec2* sample) + { + const float phi = (*sample)[1] * Pi; + const float theta = (*sample)[0] * TwoPi; + + glm::vec3 result; + const float sinPhi = sin(phi); + result[0] = cos(theta) * sinPhi; + result[1] = cos(phi); + result[2] = sin(theta) * sinPhi; + return result; + } + + __host__ __device__ float SquareToSphereUniformPDF(const glm::vec3* sample) + { + return (Inv4Pi); + } + + __host__ __device__ glm::vec3 SquareToSphereCapUniform(const glm::vec2* sample, float thetaMin) + { + const float phi = (*sample)[1] * thetaMin * Pi / 180.f; + const float theta = (*sample)[0] * 2.f * Pi; + + glm::vec3 result; + result[0] = cos(theta) * sin(phi); + result[1] = cos(phi); + result[2] = sin(theta) * sin(phi); + return result; + } + + __host__ __device__ float SquareToSphereCapUniformPDF(const glm::vec2* sample, float thetaMin) + { + return (Inv2Pi) / (1.f - cos(glm::radians(180.f - thetaMin))); + } + + __host__ __device__ glm::vec3 SquareToHemisphereUniform(const glm::vec2* sample) + { + const float z = (*sample)[0]; + const float r = glm::sqrt(glm::max(0.f, 1.f - z * z)); + const float phi = 2 * 3.14159265 * (*sample)[1]; + return glm::vec3(r * std::cos(phi), r * std::sin(phi), z); + } + + __host__ __device__ float SquareToHemisphereUniformPDF(const glm::vec2* sample) + { + return (Inv2Pi); + } + + __host__ __device__ glm::vec3 SquareToHemisphereCosine(const glm::vec2* sample) + { + const glm::vec3 d = SquareToDiskConcentric(sample); + const float z = glm::sqrt(glm::max(0.f, 1.f - d.x * d.x - d.y * d.y)); + return glm::vec3(d.x, d.y, z); + } + + __host__ __device__ float SquareToHemisphereCosinePDF(const glm::vec2* sample) + { + return (Inv8Pi); + } + +} // namespace WarpFunctions end + + + + + + + diff --git a/src/warpfunctions.h b/src/warpfunctions.h new file mode 100644 index 0000000..1ceaeba --- /dev/null +++ b/src/warpfunctions.h @@ -0,0 +1,35 @@ +#pragma once + +namespace WarpFunctions +{ + + glm::vec3 SquareToDiskUniform(const glm::vec2* sample); + + glm::vec3 SquareToDiskConcentric(const glm::vec2* sample); + + float SquareToDiskPDF(const glm::vec3* sample); + + glm::vec3 SquareToSphereUniform(const glm::vec2* sample); + + float SquareToSphereUniformPDF(const glm::vec3* sample); + + glm::vec3 SquareToSphereCapUniform(const glm::vec2* sample, float thetaMin); + + float SquareToSphereCapUniformPDF(const glm::vec2* sample, float thetaMin); + + glm::vec3 SquareToHemisphereUniform(const glm::vec2* sample); + + float SquareToHemisphereUniformPDF(const glm::vec2* sample); + + glm::vec3 SquareToHemisphereCosine(const glm::vec2* sample); + + float SquareToHemisphereCosinePDF(const glm::vec2* sample); + +} // namespace WarpFunctions end + + + + + + + diff --git a/stream_compaction/CMakeLists.txt b/stream_compaction/CMakeLists.txt index ac358c9..3a3d54c 100644 --- a/stream_compaction/CMakeLists.txt +++ b/stream_compaction/CMakeLists.txt @@ -3,5 +3,5 @@ set(SOURCE_FILES cuda_add_library(stream_compaction ${SOURCE_FILES} - OPTIONS -arch=sm_20 + OPTIONS -arch=sm_60 )