diff --git a/README.md b/README.md index 110697c..045b46a 100644 --- a/README.md +++ b/README.md @@ -1,13 +1,93 @@ -CUDA Path Tracer -================ -**University of Pennsylvania, CIS 565: GPU Programming and Architecture, Project 3** +CUDA PATH TRACER +================================================================== -* (TODO) YOUR NAME HERE -* Tested on: (TODO) Windows 22, i7-2222 @ 2.22GHz 22GB, GTX 222 222MB (Moore 2222 Lab) +**University of Pennsylvania, CIS 565: GPU Programming and Architecture** -### (TODO: Your README) +Dhruv Karthik: [LinkedIn](https://www.linkedin.com/in/dhruv_karthik/) -*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. +Tested on: Windows 10 Home, Intel(R) Core(TM) i7-8700 CPU @ 3.20GHz, 16GM, GTX 2070 - Compute Capability 7.5 +____________________________________________________________________________________ +![Developer](https://img.shields.io/badge/Developer-Dhruv-0f97ff.svg?style=flat) ![CUDA 10.1](https://img.shields.io/badge/CUDA-10.1-yellow.svg) ![Built](https://img.shields.io/appveyor/ci/gruntjs/grunt.svg) ![Issues](https://img.shields.io/badge/issues-none-green.svg) +____________________________________________________________________________________ +

+ +

+ +Table of contents +================= + * [What is Path Tracing](#what-is-path-tracing) + * [Features Overview](#features-overview) + * [BSDF Scattering: Diffuse, Specular-Reflective, Specular Transmissive](#bsdf-scattering) + * [Procedural Shapes](#procedural-shapes) + * [Stochastic Sampled Anti Aliasing](#stochastic-sampled-anti-aliasing) + * [Motion Blur](#motion-blur) + * [Optimizations ](#optimizations) + * [Stream compaction to remove terminated rays](#stream-compaction-to-remove-terminated-rays) + * [First bounce caching](#first-bounce-caching) + * [Sort by Material](#sort-by-material) + * [Questions](#questions) + * [Performance Analysis](#performance-analysis) + * [Credits & Acknowledgments](#credits) + +# What is path tracing? +Path tracing refers to a set of techniques to virtually render images by emulating certain physical properties of light. In real life, Rays of light leave light sources, bounce around the world, and hit pixels in the camera. Path traces simulates this effect by firing 'rays' out of the camera pixels, and considering those that hit a light source. +

+Pathtrace +

+ +# Features Overview +## BSDF Scattering +A combination of ***reflection*** and ***transmission*** functions that describe how rays must bounce once they intersect an object. For transmissive and refractive objects, I used schlicks approximation to calculate the probability of the refractive surface being reflective at high incidence angles. The following illustrates BSDF on three material properties: + +| Diffuse | Reflective | Refractive | +| ------------- | ----------- |----------- | +| ![](img/bsdf3.png) | ![](img/bsdf1.png) | ![](img/bsdf2.png) | + +## Procedural Shapes +I created procedural shapes via a variation of ***Constructive Solid Geometry*** ray tracing. Using the provided primitives, I modified in the code in ```intersection.h``` to generate three different constructive of the original primitive geometry that used portions of the original geometry and negated others: + +| 1 | 2 | 3 | +| ------------- | ----------- |----------- | +| ![](img/SphereAndNotCube.PNG) | ![](img/SphereCubeUnion.PNG) | ![](img/SphereAndCubeMatrix.PNG) | + +## Stochastic Sampled Anti Aliasing +Implementing this simply required that randomly I nudge the ray direction by epsilon in the x & and y axis. Notice how the pre-AA image has jagged edges where the yellow light reflects off it. Notice how this issue dissappears post AA! + +| Pre AA | Post AA | +| ------------- | ----------- | +| ![](img/preAA.jpg) | ![](img/postAA.jpg) | + +## Motion Blur +I created a kernel function ```blurGeom``` that allows users to specify a ``glm::vec3 offset`` so they could direct the blur towards their desired position. I implemented motion blur by targeting a subset of ``geom`` objects to respond to changes in ``dt = epsilon * iter``. The second run illustrates motion blur with first bounce cache on - This is not the right thing to do, but resulted in a cool image. + +| 1 | 2 | 3 | +| ------------- | ----------- | ----------- | +| ![](img/mblurmirror.PNG) | ![](img/mirrordim.PNG) | ![](img/motionblur.PNG) | + +# Optimizations +## Stream Compaction to Remove Terminated Rays +CUDA can only launch a fininite number of blocks at a time. As a result, if some threads are tracing more bounces, and some are tracing a few, we can end up with a lot of idling threads. One solution to this is to stream compact away idling threads. This means terminated threads are removed at the every kernel's launch, so that running threads are grouped together. +## First Bounce Caching +This is a clever optimization to avoid recomputing the first bounce for depth=0. This is because this bounce stays the same regardless of iterations, so we precompute this in the first bounce and reload it whenvever depth=0. +## Sort By Material +Divergence is a seriously perfomance killer in CUDA, as if different threads in a warp diverge then the benefits of parralelilzation begin to dissapear. This would most likely happen in a path tracing scenario when materials that reflect differently run at the same time. To avoid warp divergence while the kernel is running, I sorted by material to ensure that similar materials most likley get placed in the same warp as they are contiguous in memory. + +# Performance Analysis +I proceeded to remove walls from my cornell scene and to printed out the decrease in number of elements left in the array post stream compaction. Removing walls exposed more black space and thereby would result in a more rapid decrease of elements in the compacted array. I started with only 2 walls, and this had the greatest decrease as it started at approx 2000 values and went to near zero in about 4 bounces. + +# Questions + +**Compare scenes which are open (like the given cornell box) and closed (i.e. no light can escape the scene). Again, compare the performance effects of stream compaction! Remember, stream compaction only affects rays which terminate, so what might you expect?** +If no light can escape the scene, then the benefits of stream compaction are nullified as rays only terminate when they reach the end of their depth. As a result, the presence of stream compaction had no effect on the closed box scene. + +# Bloopers +See wayyy more images and bloopers in the imgs folder! +Messed up refraction +![](img/something.PNG) + + +# Credits +[rayttracing.github.io](raytracing.github.io) diff --git a/img/0.PNG b/img/0.PNG new file mode 100644 index 0000000..618a4fd Binary files /dev/null and b/img/0.PNG differ diff --git a/img/1.PNG b/img/1.PNG new file mode 100644 index 0000000..98d46c5 Binary files /dev/null and b/img/1.PNG differ diff --git a/img/2.PNG b/img/2.PNG new file mode 100644 index 0000000..64755d2 Binary files /dev/null and b/img/2.PNG differ diff --git a/img/3.PNG b/img/3.PNG new file mode 100644 index 0000000..8ea1c70 Binary files /dev/null and b/img/3.PNG differ diff --git a/img/Antialiasing.PNG b/img/Antialiasing.PNG new file mode 100644 index 0000000..5340442 Binary files /dev/null and b/img/Antialiasing.PNG differ diff --git a/img/MirrorDimension.PNG b/img/MirrorDimension.PNG new file mode 100644 index 0000000..e92be52 Binary files /dev/null and b/img/MirrorDimension.PNG differ diff --git a/img/Motion Blur Mirror Dimension2.PNG b/img/Motion Blur Mirror Dimension2.PNG new file mode 100644 index 0000000..6f75857 Binary files /dev/null and b/img/Motion Blur Mirror Dimension2.PNG differ diff --git a/img/MotionBlurMirrorDim.PNG b/img/MotionBlurMirrorDim.PNG new file mode 100644 index 0000000..165c295 Binary files /dev/null and b/img/MotionBlurMirrorDim.PNG differ diff --git a/img/Sphere and NotCube.PNG b/img/Sphere and NotCube.PNG new file mode 100644 index 0000000..ff21332 Binary files /dev/null and b/img/Sphere and NotCube.PNG differ diff --git a/img/SphereAndCubeMatrix.PNG b/img/SphereAndCubeMatrix.PNG new file mode 100644 index 0000000..82c915f Binary files /dev/null and b/img/SphereAndCubeMatrix.PNG differ diff --git a/img/SphereAndNotCube.PNG b/img/SphereAndNotCube.PNG new file mode 100644 index 0000000..3141eca Binary files /dev/null and b/img/SphereAndNotCube.PNG differ diff --git a/img/SphereCubeUnion.PNG b/img/SphereCubeUnion.PNG new file mode 100644 index 0000000..7f18a2f Binary files /dev/null and b/img/SphereCubeUnion.PNG differ diff --git a/img/SphereCubeUnion2.PNG b/img/SphereCubeUnion2.PNG new file mode 100644 index 0000000..8802072 Binary files /dev/null and b/img/SphereCubeUnion2.PNG differ diff --git a/img/Spotlight Image.PNG b/img/Spotlight Image.PNG new file mode 100644 index 0000000..0363f98 Binary files /dev/null and b/img/Spotlight Image.PNG differ diff --git a/img/bsdf1.png b/img/bsdf1.png new file mode 100644 index 0000000..01e3b9c Binary files /dev/null and b/img/bsdf1.png differ diff --git a/img/bsdf2.png b/img/bsdf2.png new file mode 100644 index 0000000..7935a75 Binary files /dev/null and b/img/bsdf2.png differ diff --git a/img/bsdf3.png b/img/bsdf3.png new file mode 100644 index 0000000..4607c6a Binary files /dev/null and b/img/bsdf3.png differ diff --git a/img/frontpage.png b/img/frontpage.png new file mode 100644 index 0000000..86a9c89 Binary files /dev/null and b/img/frontpage.png differ diff --git a/img/mblurmirror.PNG b/img/mblurmirror.PNG new file mode 100644 index 0000000..d87f77e Binary files /dev/null and b/img/mblurmirror.PNG differ diff --git a/img/mirrordim.PNG b/img/mirrordim.PNG new file mode 100644 index 0000000..2261b77 Binary files /dev/null and b/img/mirrordim.PNG differ diff --git a/img/motionblur.PNG b/img/motionblur.PNG new file mode 100644 index 0000000..8c3ad58 Binary files /dev/null and b/img/motionblur.PNG differ diff --git a/img/postAA.jpg b/img/postAA.jpg new file mode 100644 index 0000000..c07f1f8 Binary files /dev/null and b/img/postAA.jpg differ diff --git a/img/pre-aa.png b/img/pre-aa.png new file mode 100644 index 0000000..f773a1e Binary files /dev/null and b/img/pre-aa.png differ diff --git a/img/preAA.jpg b/img/preAA.jpg new file mode 100644 index 0000000..ff6c1e4 Binary files /dev/null and b/img/preAA.jpg differ diff --git a/img/something.PNG b/img/something.PNG new file mode 100644 index 0000000..65542bf Binary files /dev/null and b/img/something.PNG differ diff --git a/img/wasted.PNG b/img/wasted.PNG new file mode 100644 index 0000000..03c46c2 Binary files /dev/null and b/img/wasted.PNG differ diff --git a/scenes/cornell.txt b/scenes/cornell.txt index 83ff820..70f073e 100644 --- a/scenes/cornell.txt +++ b/scenes/cornell.txt @@ -6,7 +6,7 @@ SPECRGB 0 0 0 REFL 0 REFR 0 REFRIOR 0 -EMITTANCE 5 +EMITTANCE 10 // Diffuse white MATERIAL 1 @@ -48,6 +48,18 @@ REFR 0 REFRIOR 0 EMITTANCE 0 +// Glass white +MATERIAL 5 +RGB 1 1 1 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0.5 +REFR 0.5 +REFRIOR 1.50s +EMITTANCE 0 + + + // Camera CAMERA RES 800 800 @@ -95,7 +107,7 @@ SCALE .01 10 10 // Left wall OBJECT 4 cube -material 2 +material 1 TRANS -5 5 0 ROTAT 0 0 0 SCALE .01 10 10 @@ -111,7 +123,7 @@ SCALE .01 10 10 // Sphere OBJECT 6 sphere -material 4 -TRANS -1 4 -1 +material 5 +TRANS -1 5 -1 ROTAT 0 0 0 -SCALE 3 3 3 +SCALE 5 5 5 \ No newline at end of file diff --git a/scenes/fourlightroom.txt b/scenes/fourlightroom.txt new file mode 100644 index 0000000..224b1d1 --- /dev/null +++ b/scenes/fourlightroom.txt @@ -0,0 +1,219 @@ +// Emissive Blue material (light) +MATERIAL 0 +RGB 0.8 0.8 0.0 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 10 + +// Emissive Red material (light) +MATERIAL 1 +RGB 1 0.0 0.3 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 5 + + +// Emissive Blue material (light) +MATERIAL 2 +RGB 0 1 0.5 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 5 + +// Emissive Green material (light) +MATERIAL 3 +RGB 0.01 0.95 0.1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 10 + +// Diffuse white +MATERIAL 4 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Diffuse red +MATERIAL 5 +RGB .85 .35 .35 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Diffuse green +MATERIAL 6 +RGB .1 1.0 0.1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Diffuse Dark Blue +MATERIAL 7 +RGB 0.9 0.5 0.9 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Reflection & Refraction +MATERIAL 8 +RGB 0.8 0.4 0.4 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0.7 +REFR 0.3 +REFRIOR 1.35 +EMITTANCE 0 + +// Reflection & Refraction +MATERIAL 9 +RGB 0.8 0.6 0.8 +SPECEX 0 +SPECRGB 1.0 1.0 1.0 +REFL 1.0 +REFR 0.0 +REFRIOR 0 +EMITTANCE 0 + +// Reflection & Refraction +MATERIAL 10 +RGB 1.0 1.0 1.0 +SPECEX 0 +SPECRGB 1.0 1.0 1.0 +REFL 0.5 +REFR 0.5 +REFRIOR 1.56 +EMITTANCE 0 + +// Reflection & Refraction +MATERIAL 11 +RGB 0.1 1.0 1.0 +SPECEX 0 +SPECRGB 1.0 1.0 1.0 +REFL 0.9 +REFR 0.1 +REFRIOR 1.2 +EMITTANCE 0 + + +// Camera +CAMERA +RES 800 800 +FOVY 45 +ITERATIONS 5000 +DEPTH 8 +FILE fourlightroom +EYE 0.0 5 10.5 +LOOKAT 0 5 0 +UP 0 1 0 + + +// Ceiling Blue Light +OBJECT 0 +cube +material 0 +TRANS 0 10 0 +ROTAT 0 0 0 +SCALE 3 .3 3 + + +// Left Pink Light +OBJECT 1 +cube +material 1 +TRANS -5 500 0 +ROTAT 0 0 90 +SCALE 3 .3 3 + + +// Right Yellow light +OBJECT 2 +cube +material 2 +TRANS 5 500 0 +ROTAT 0 0 90 +SCALE 3 .3 3 + + +// Floor Green light +OBJECT 3 +cube +material 3 +TRANS 0 0 0 +ROTAT 0 0 0 +SCALE 3 .3 3 + + +// Floor +OBJECT 4 +cube +material 7 +TRANS 0 0 0 +ROTAT 0 0 180 +SCALE 10 .01 10 + +// Ceiling +OBJECT 5 +cube +material 7 +TRANS 0 10 0 +ROTAT 0 0 90 +SCALE .01 10 10 + +// Back wall +OBJECT 6 +cube +material 7 +TRANS 0 5 -5 +ROTAT 0 90 0 +SCALE .01 10 10 + +// Left wall +OBJECT 7 +cube +material 9 +TRANS -5 5 0 +ROTAT 0 0 0s +SCALE .01 10 10 + +// Right wall +OBJECT 8 +cube +material 9 +TRANS 5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + + +// SphereandNotCUBE +OBJECT 9 +sphereAndNotCube +material 8 +TRANS -1 4 -1 +ROTAT 0 10 45 +SCALE 3 3 3 + diff --git a/scenes/frontpage.txt b/scenes/frontpage.txt new file mode 100644 index 0000000..3f45428 --- /dev/null +++ b/scenes/frontpage.txt @@ -0,0 +1,107 @@ +// 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 + +// Glass white +MATERIAL 5 +RGB 1 1 1 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 1 +REFRIOR 1.55 +EMITTANCE 0 + +// Floor +OBJECT 1 +cube +material 1 +TRANS 0 0 0 +ROTAT 0 0 0 +SCALE 10 .01 10 + +// Ceiling +OBJECT 2 +cube +material 1 +TRANS 0 10 0 +ROTAT 0 0 90 +SCALE .01 10 10 + +// Back wall +OBJECT 3 +cube +material 1 +TRANS 0 5 -5 +ROTAT 0 90 0 +SCALE .01 10 10 + +// Left wall +OBJECT 4 +cube +material 2 +TRANS -5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +// Right wall +OBJECT 5 +cube +material 3 +TRANS 5 5%% 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +// Sphere +OBJECT 6 +sphere +material 4 +TRANS -1 4 -1 +ROTAT 0 0 0 +SCALE 3 3 3 \ No newline at end of file diff --git a/scenes/refractTest.txt b/scenes/refractTest.txt new file mode 100644 index 0000000..c8fe413 --- /dev/null +++ b/scenes/refractTest.txt @@ -0,0 +1,127 @@ +// 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 + +// Glass +MATERIAL 4 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB .98 .98 .98 +REFL 0 +REFR 1 +REFRIOR 1.45 +EMITTANCE 0 + +// Water +MATERIAL 5 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB .98 .98 .98 +REFL 0 +REFR 1 +REFRIOR 1.35 +EMITTANCE 0 + +// Camera +CAMERA +RES 800 800 +FOVY 45 +ITERATIONS 5000 +DEPTH 8 +FILE cornell +EYE 0.0 5 10.5 +LOOKAT 0 5 0 +UP 0 1 0 + + +// Ceiling light +OBJECT 0 +cube +material 0 +TRANS 0 10 0 +ROTAT 0 0 0 +SCALE 3 .3 3 + +// Floor +OBJECT 1 +cube +material 1 +TRANS 0 0 0 +ROTAT 0 0 0 +SCALE 10 .01 10 + +// Ceiling +OBJECT 2 +cube +material 1 +TRANS 0 10 0 +ROTAT 0 0 90 +SCALE .01 10 10 + +// Back wall +OBJECT 3 +cube +material 1 +TRANS 0 5 -5 +ROTAT 0 90 0 +SCALE .01 10 10 + +// Left wall +OBJECT 4 +cube +material 2 +TRANS -5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +// Right wall +OBJECT 5 +cube +material 3 +TRANS 5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +// Sphere +OBJECT 6 +sphere +material 5 +TRANS -1 4 -1 +ROTAT 0 0 0 +SCALE 3 3 3 diff --git a/src/interactions.h b/src/interactions.h index 5ce3628..15cfd25 100644 --- a/src/interactions.h +++ b/src/interactions.h @@ -66,6 +66,13 @@ glm::vec3 calculateRandomDirectionInHemisphere( * * You may need to change the parameter list for your purposes! */ + + +__host__ __device__ +float schlick(float cosine, float ref_idx) { + float r0 = powf((1.f - ref_idx) / (1.f + ref_idx), 2.f); + return r0 + (1.f - r0) * powf((1.f - cosine), 5.f); +} __host__ __device__ void scatterRay( PathSegment & pathSegment, @@ -73,7 +80,47 @@ void scatterRay( glm::vec3 normal, const Material &m, thrust::default_random_engine &rng) { - // TODO: implement this. - // A basic implementation of pure-diffuse shading will just call the - // calculateRandomDirectionInHemisphere defined above. -} + + thrust::uniform_real_distribution u01(0, 1); + float prob = u01(rng); + if (prob < m.hasRefractive) { + //Refract + glm::vec3 inDirection = pathSegment.ray.direction; + //Dot is positive if normal & inDirection face the same dir -> the ray is inside the object getting out + bool insideObj = glm::dot(inDirection, normal) > 0.0f; + + //glm::refract (followed raytracing.github.io trick for hollow glass sphere effect by reversing normals) + float eta = insideObj ? m.indexOfRefraction : (1.0f / m.indexOfRefraction); + glm::vec3 outwardNormal = insideObj ? -1.0f * normal : normal; + glm::vec3 finalDir = glm::refract(glm::normalize(inDirection), glm::normalize(outwardNormal), eta); + + //Check for TIR (if magnitude of refracted ray is very small) + if (glm::length(finalDir) < 0.01f) { + pathSegment.color *= 0.0f; + finalDir = glm::reflect(inDirection, normal); + } + + //Use schlicks to calculate reflective probability (also followed raytracing.github.io) + float cosine = glm::dot(inDirection, normal); + float reflectProb = schlick(cosine, m.indexOfRefraction); + float sampleFloat = u01(rng); + + pathSegment.ray.direction = reflectProb < sampleFloat ? glm::reflect(inDirection, normal) : finalDir; + pathSegment.ray.origin = intersect + 0.001f * pathSegment.ray.direction; + pathSegment.color *= m.specular.color; + } + else if (prob < m.hasReflective) { + //Reflective Surface + pathSegment.ray.direction = glm::reflect(pathSegment.ray.direction, normal); + pathSegment.ray.origin = intersect + 0.01f * normal; + pathSegment.color *= m.specular.color; + } + else { + //Diffuse Surface + pathSegment.ray.direction = calculateRandomDirectionInHemisphere(normal, rng); + pathSegment.ray.origin = intersect + EPSILON * normal; + } + pathSegment.remainingBounces--; + pathSegment.color *= m.color; + pathSegment.color = glm::clamp(pathSegment.color, glm::vec3(0.0f), glm::vec3(1.0f)); +} \ No newline at end of file diff --git a/src/intersections.h b/src/intersections.h index 6f23872..07a8771 100644 --- a/src/intersections.h +++ b/src/intersections.h @@ -101,7 +101,8 @@ __host__ __device__ float boxIntersectionTest(Geom box, Ray r, */ __host__ __device__ float sphereIntersectionTest(Geom sphere, Ray r, glm::vec3 &intersectionPoint, glm::vec3 &normal, bool &outside) { - float radius = .5; + //float radius = .5; + float radius = 0.55; 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))); @@ -111,6 +112,9 @@ __host__ __device__ float sphereIntersectionTest(Geom sphere, Ray r, rt.direction = rd; float vDotDirection = glm::dot(rt.origin, rt.direction); + //glm::dot(rt.origin, rt.origin): dot(A-C, A-C) + //powf(radius, 2) = R^2 + //vDotDirection = dot(B, A-C) float radicand = vDotDirection * vDotDirection - (glm::dot(rt.origin, rt.origin) - powf(radius, 2)); if (radicand < 0) { return -1; @@ -142,3 +146,121 @@ __host__ __device__ float sphereIntersectionTest(Geom sphere, Ray r, return glm::length(r.origin - intersectionPoint); } + + +__host__ __device__ float sphereIntersectionTestRadius(Geom sphere, Ray r, + glm::vec3 &intersectionPoint, glm::vec3 &normal, bool &outside, float radius) { + + 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); + //glm::dot(rt.origin, rt.origin): dot(A-C, A-C) + //powf(radius, 2) = R^2 + //vDotDirection = dot(B, A-C) + 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); +} + +// CHECKITOUT +/** + * Test intersection between a ray and Union of Sphere & Cube. Untransformed, + * the sphere always has radius 0.5, and untransformed cube has side lenghts 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 sphereCubeUnion(Geom spherecube, Ray r, + glm::vec3 &intersectionPoint, glm::vec3 &normal, bool &outside) { + glm::vec3 cubeIntersection, cubeNormal; + glm::vec3 sphereIntersection, sphereNormal; + bool cubeOutside, sphereOutside; + float tCube, tSphere; + tSphere = sphereIntersectionTestRadius(spherecube, r, sphereIntersection, sphereNormal, sphereOutside, 0.65); + tCube = boxIntersectionTest(spherecube, r, cubeIntersection, cubeNormal, cubeOutside); + if (tSphere <= tCube) { + intersectionPoint = sphereIntersection; + normal = sphereNormal; + outside = sphereOutside; + return tSphere; + } + else { + intersectionPoint = cubeIntersection; + normal = cubeNormal; + outside = cubeOutside; + return tCube; + } +} + + +__host__ __device__ float sphereAndCube(Geom spherecube, Ray r, + glm::vec3 &intersectionPoint, glm::vec3 &normal, bool &outside) { + glm::vec3 cubeIntersection, cubeNormal; + glm::vec3 sphereIntersection, sphereNormal; + bool cubeOutside, sphereOutside; + float tCube, tSphere; + tSphere = sphereIntersectionTestRadius(spherecube, r, sphereIntersection, sphereNormal, sphereOutside, 0.55); + tCube = boxIntersectionTest(spherecube, r, cubeIntersection, cubeNormal, cubeOutside); + if (tCube <= tSphere) { + intersectionPoint = sphereIntersection; + normal = sphereNormal; + outside = sphereOutside; + return tSphere; + } + else { + return -1; + } +} + +__host__ __device__ float sphereAndNotCube(Geom spherecube, Ray r, + glm::vec3 &intersectionPoint, glm::vec3 &normal, bool &outside) { + glm::vec3 cubeIntersection, cubeNormal; + glm::vec3 sphereIntersection, sphereNormal; + bool cubeOutside, sphereOutside; + float tCube, tSphere; + tSphere = sphereIntersectionTestRadius(spherecube, r, sphereIntersection, sphereNormal, sphereOutside, 0.60); + tCube = boxIntersectionTest(spherecube, r, cubeIntersection, cubeNormal, cubeOutside); + if (tSphere <= tCube) { + intersectionPoint = sphereIntersection; + normal = sphereNormal; + outside = sphereOutside; + return tSphere; + } + else { + return -1; + } +} diff --git a/src/pathtrace.cu b/src/pathtrace.cu index c1ec122..55414bd 100644 --- a/src/pathtrace.cu +++ b/src/pathtrace.cu @@ -1,20 +1,31 @@ #include #include #include +#include +#include #include #include +#include #include #include "sceneStructs.h" #include "scene.h" #include "glm/glm.hpp" +#include "glm/gtc/matrix_transform.hpp" #include "glm/gtx/norm.hpp" +#include "glm/matrix.hpp" +#include "glm/glm.hpp" #include "utilities.h" #include "pathtrace.h" #include "intersections.h" #include "interactions.h" #define ERRORCHECK 1 +#define ANTIALIASING 0 +#define SORTMATERIAL 1 +#define STREAMCOMPACT 1 +#define BLURGEOM 0 +#define CACHE 1 #define FILENAME (strrchr(__FILE__, '/') ? strrchr(__FILE__, '/') + 1 : __FILE__) #define checkCUDAError(msg) checkCUDAErrorFn(msg, FILENAME, __LINE__) @@ -73,8 +84,8 @@ static Geom * dev_geoms = NULL; static Material * dev_materials = NULL; static PathSegment * dev_paths = NULL; static ShadeableIntersection * dev_intersections = NULL; +static ShadeableIntersection * dev_intersections_cache = NULL; // TODO: static variables for device memory, any extra info you need, etc -// ... void pathtraceInit(Scene *scene) { hst_scene = scene; @@ -95,6 +106,8 @@ void pathtraceInit(Scene *scene) { cudaMalloc(&dev_intersections, pixelcount * sizeof(ShadeableIntersection)); cudaMemset(dev_intersections, 0, pixelcount * sizeof(ShadeableIntersection)); + cudaMalloc(&dev_intersections_cache, pixelcount * sizeof(ShadeableIntersection)); + cudaMemset(dev_intersections_cache, 0, pixelcount * sizeof(ShadeableIntersection)); // TODO: initialize any extra device memeory you need checkCUDAError("pathtraceInit"); @@ -106,6 +119,7 @@ void pathtraceFree() { cudaFree(dev_geoms); cudaFree(dev_materials); cudaFree(dev_intersections); + cudaFree(dev_intersections_cache); // TODO: clean up any extra device memory you created checkCUDAError("pathtraceFree"); @@ -129,19 +143,41 @@ __global__ void generateRayFromCamera(Camera cam, int iter, int traceDepth, Path 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); + float deltax, deltay; + thrust::default_random_engine rng = makeSeededRandomEngine(iter, index, segment.remainingBounces); + thrust::uniform_real_distribution u01(0, 1); +#if ANTIALIASING + deltax = u01(rng); + deltay = u01(rng); +#endif // 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 * ((float)x - (float)cam.resolution.x * 0.5f + deltax) + - cam.up * cam.pixelLength.y * ((float)y - (float)cam.resolution.y * 0.5f + deltay) ); - segment.pixelIndex = index; segment.remainingBounces = traceDepth; } } +__global__ void blurGeom(Geom *geoms, int geoms_size, int num_paths, glm::vec3 offset, int iter) { + int path_index = blockIdx.x * blockDim.x + threadIdx.x; + if (path_index < num_paths) { + float dt = iter * 0.00001f; + for (int i = 0; i < geoms_size; ++i) { + Geom & geom = geoms[i]; + if (geom.type == SPHERE || geom.type == SPHEREANDCUBE) { + geom.translation -= glm::clamp(offset * dt, glm::vec3(0.0f), offset); + geom.transform[3] = glm::vec4(geom.translation, geom.transform[3].w); + geom.inverseTransform = glm::inverse(geom.transform); + geom.invTranspose = glm::transpose(geom.inverseTransform); + } + } + } +} + // TODO: // computeIntersections handles generating ray intersections ONLY. // Generating new rays is handled in your shader(s). @@ -152,7 +188,7 @@ __global__ void computeIntersections( , PathSegment * pathSegments , Geom * geoms , int geoms_size - , ShadeableIntersection * intersections + , ShadeableIntersection * intersections ) { int path_index = blockIdx.x * blockDim.x + threadIdx.x; @@ -176,15 +212,22 @@ __global__ void computeIntersections( for (int i = 0; i < geoms_size; i++) { Geom & geom = geoms[i]; - if (geom.type == CUBE) { t = boxIntersectionTest(geom, pathSegment.ray, tmp_intersect, tmp_normal, outside); } - else if (geom.type == SPHERE) - { + else if (geom.type == SPHERE) { t = sphereIntersectionTest(geom, pathSegment.ray, tmp_intersect, tmp_normal, outside); } + else if (geom.type == SPHERECUBEUNION) { + t = sphereCubeUnion(geom, pathSegment.ray, tmp_intersect, tmp_normal, outside); + } + else if (geom.type == SPHEREANDNOTCUBE) { + t = sphereAndNotCube(geom, pathSegment.ray, tmp_intersect, tmp_normal, outside); + } + else if (geom.type == SPHEREANDCUBE) { + t = sphereAndCube(geom, pathSegment.ray, tmp_intersect, tmp_normal, outside); + } // TODO: add more intersection tests here... triangle? metaball? CSG? // Compute the minimum t from the intersection tests to determine what @@ -230,14 +273,14 @@ __global__ void shadeFakeMaterial ( ) { int idx = blockIdx.x * blockDim.x + threadIdx.x; - if (idx < num_paths) + if (idx < num_paths && pathSegments[idx].remainingBounces) { ShadeableIntersection intersection = shadeableIntersections[idx]; if (intersection.t > 0.0f) { // if the intersection exists... // Set up the RNG // LOOK: this is how you use thrust's RNG! Please look at // makeSeededRandomEngine as well. - thrust::default_random_engine rng = makeSeededRandomEngine(iter, idx, 0); + thrust::default_random_engine rng = makeSeededRandomEngine(iter, idx, pathSegments[idx].remainingBounces); thrust::uniform_real_distribution u01(0, 1); Material material = materials[intersection.materialId]; @@ -246,21 +289,25 @@ __global__ void shadeFakeMaterial ( // If the material indicates that the object was a light, "light" the ray if (material.emittance > 0.0f) { pathSegments[idx].color *= (materialColor * material.emittance); + pathSegments[idx].remainingBounces = 0; } // Otherwise, do some pseudo-lighting computation. This is actually more // 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)); - 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 + glm::vec3 rayObjectIntersection = getPointOnRay(pathSegments[idx].ray, intersection.t); + scatterRay(pathSegments[idx], rayObjectIntersection, intersection.surfaceNormal, material, rng); + //float lightTerm = glm::dot(intersection.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 } // If there was no intersection, color the ray black. // Lots of renderers use 4 channel color, RGBA, where A = alpha, often // used for opacity, in which case they can indicate "no opacity". // This can be useful for post-processing and image compositing. } else { - pathSegments[idx].color = glm::vec3(0.0f); + pathSegments[idx].color = glm::vec3(0.0f); + pathSegments[idx].remainingBounces = 0; } } } @@ -277,6 +324,18 @@ __global__ void finalGather(int nPaths, glm::vec3 * image, PathSegment * iterati } } +struct compareMaterials { + __host__ __device__ bool operator()(const ShadeableIntersection &m1, const ShadeableIntersection &m2) { + return m1.materialId < m2.materialId; + } +}; + +struct isPathContinue { + __host__ __device__ bool operator()(const PathSegment &p) { + return (p.remainingBounces > 0); + } +}; + /** * Wrapper for the __global__ call that sets up the kernel calls and does a ton * of memory management @@ -326,56 +385,86 @@ void pathtrace(uchar4 *pbo, int frame, int iter) { // TODO: perform one iteration of path tracing - 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; - - // --- PathSegment Tracing Stage --- - // Shoot ray into scene, bounce between objects, push shading chunks - - bool iterationComplete = false; - while (!iterationComplete) { + 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; + + // --- 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; + + #if BLURGEOM + blurGeom << > > (dev_geoms, hst_scene->geoms.size(), num_paths, glm::vec3(-1e-5f, -1e-5f, 0.0), iter); + #endif //BLURGEOM + + ShadeableIntersection* curr_Intersections = dev_intersections; + + if (depth == 0 && CACHE) { + if (iter == 1) { + computeIntersections <<>> ( + depth + , num_paths + , dev_paths + , dev_geoms + , hst_scene->geoms.size() + , dev_intersections + ); + cudaMemcpy(dev_intersections_cache, dev_intersections, pixelcount * sizeof(ShadeableIntersection), cudaMemcpyDeviceToDevice); + checkCUDAError("CUDA memcpy cache to intersections failed"); + } + else { + //reload + curr_Intersections = dev_intersections_cache; + } + } + else { + computeIntersections <<>> ( + depth + , num_paths + , dev_paths + , dev_geoms + , hst_scene->geoms.size() + , dev_intersections + ); + checkCUDAError("trace one bounce"); + } + cudaDeviceSynchronize(); + + #if SORTMATERIAL + thrust::sort_by_key(thrust::device, dev_intersections, dev_intersections + num_paths, dev_paths, compareMaterials()); + #endif // SORTMATERIAL + + shadeFakeMaterial<<>> ( + iter, + num_paths, + curr_Intersections, + dev_paths, + dev_materials + ); - // clean shading chunks - cudaMemset(dev_intersections, 0, pixelcount * sizeof(ShadeableIntersection)); + #if STREAMCOMPACT + PathSegment* newEnd = thrust::partition(thrust::device, dev_paths, dev_paths + num_paths, isPathContinue()); + num_paths = newEnd - dev_paths; + depth = (num_paths < 1) ? traceDepth + 1 : depth; + #endif //STREAMCOMPACT - // 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. + iterationComplete = depth > traceDepth || num_paths <= 0; + depth++; } + num_paths = dev_path_end - dev_paths; + // Assemble this iteration and apply it to the image dim3 numBlocksPixels = (pixelcount + blockSize1d - 1) / blockSize1d; finalGather<<>>(num_paths, dev_image, dev_paths); @@ -390,4 +479,4 @@ void pathtrace(uchar4 *pbo, int frame, int iter) { pixelcount * sizeof(glm::vec3), cudaMemcpyDeviceToHost); checkCUDAError("pathtrace"); -} +} \ No newline at end of file diff --git a/src/scene.cpp b/src/scene.cpp index cbae043..5fa70a2 100644 --- a/src/scene.cpp +++ b/src/scene.cpp @@ -51,6 +51,18 @@ 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(), "sphereCubeUnion") == 0) { + cout << "Creating new sphereCubeUnion..." << endl; + newGeom.type = SPHERECUBEUNION; + } + else if (strcmp(line.c_str(), "sphereAndCube") == 0) { + cout << "Creating new sphereAndCube..." << endl; + newGeom.type = SPHEREANDCUBE; + } + else if (strcmp(line.c_str(), "sphereAndNotCube") == 0) { + cout << "Creating new sphereAndNotCube..." << endl; + newGeom.type = SPHEREANDNOTCUBE; } } @@ -125,7 +137,6 @@ int Scene::loadCamera() { } else if (strcmp(tokens[0].c_str(), "UP") == 0) { camera.up = glm::vec3(atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str())); } - utilityCore::safeGetline(fp_in, line); } diff --git a/src/sceneStructs.h b/src/sceneStructs.h index b38b820..41f2e8d 100644 --- a/src/sceneStructs.h +++ b/src/sceneStructs.h @@ -10,6 +10,9 @@ enum GeomType { SPHERE, CUBE, + SPHERECUBEUNION, + SPHEREANDCUBE, + SPHEREANDNOTCUBE }; struct Ray {