diff --git a/README.md b/README.md index 110697c..fe10750 100644 --- a/README.md +++ b/README.md @@ -3,11 +3,77 @@ CUDA Path Tracer **University of Pennsylvania, CIS 565: GPU Programming and Architecture, Project 3** -* (TODO) YOUR NAME HERE -* Tested on: (TODO) Windows 22, i7-2222 @ 2.22GHz 22GB, GTX 222 222MB (Moore 2222 Lab) +* Eyad Almoamen + * [LinkedIn](https://www.linkedin.com/in/eyadalmoamen/), [personal website](https://eyadnabeel.com) +* Tested on: Windows 11, i7-10750H CPU @ 2.60GHz 2.59 GHz 16GB, RTX 2070 Super Max-Q Design 8GB (Personal Computer) -### (TODO: Your README) +Introduction +================ +I've built a GPU accelerated monte carlo path tracer using CUDA and C++. The parallelization is happening on a ray-by-ray basis, with the terminated rays being eliminated via stream compaction and sorted by material type in order to avoid warp divergence. The path tracer takes in a scene description .txt file and outputs a rendered image. + +Features implemented include: + +* [Specular Reflective Material](#specular-reflective-material) +* [Refractive Material](#refractive-material) +* [Thin Lens Model DOF](#thin-lens-model-dof) +* [Motion Blur](#motion-blur) +* [Stochastic Antialiasing](#stochastic-antialiasing) +* [Direct Lighting](#direct-lighting) + +#Specular Reflective Material +================ +The specular reflective material either reflects light perfectly (incident angle == exitent angle), or diffusely, the rate of each is manually set and the two percentages sum up to 100% (for example, if the material was 63% specular, it'd have to be 37% diffuse): + + + +#Refractive Material +================ +The specular refractive material either reflects light or transmits it according to [Snell's Law](https://en.wikipedia.org/wiki/Snell%27s_law), the rate of each is based on the material type and index of refration. This is usually calculated by the [Fresnel Equations](https://en.wikipedia.org/wiki/Fresnel_equations), however, here I use the [Schlick approximation](https://en.wikipedia.org/wiki/Schlick%27s_approximation) to calculate the rates as it's more computationally efficient with a very low error rate: + + + + + +#Thin Lens Model DOF +================ +I utilized the [Thin Lens Model](https://pbr-book.org/3ed-2018/Camera_Models/Projective_Camera_Models#TheThinLensModelandDepthofField) in order to replace the pinhole camera we have with a more realistic virtual lens which allows me to introduce depth of field effects and bokeh: + +| Focal Distance | 0 | 3 | 8.5 | 20.5 | +| :------- | :-------: | :-------: | :-------: | :-------: | +| Iterations | 7759 | 5082 | 5142 | 5009 | +| Scene | | | | | + +#Motion Blur +================ +I added a velocity component to the geometry struct and that allows me to render the image in such a way that it seems the object is moving in the direction of the velocity: + +#Stochastic Antialiasing +================ +I added support for stochastic antialiasing by jittering the ray produced from the camera randomly within the range of a pixel length: + +| Antialiasing | Without | With | +| :------- | :-------: | :-------: | +| Scene | | | +| Scene | | | + +#Direct Lighting +================ +To optimize the result and speed up the convergence of the image, I had the pathtracer trace its last ray to a light source in the scene, guaranteeing that we get light contribution. To demonstrate, I've rendered the same scene up to 1000 iterations with and without direct lighting: + +| Direct Lighting | Without | With | +| :------- | :-------: | :-------: | +| Scene | | | + +Performance Testing +================ +I ran a few tests to see the effect of some of the optimizations I've performed on this path tracer: + +The effect of caching is very much evident and it increases as the size of the image increases: + + + +This is because we're precomputing a potentially very large computation, sparing ourselves the trouble for upcoming iterations -*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. +The effect of material sorting doesn't seem to be too encouraging; initially I tried testing it on a scene with one material, it wasn't an improvement (since we'd be sorting to avoid nonexistent warp divergence). However I switched to a scene with diffuse, reflective, and refractive material to no avail: + diff --git a/img/Screenshot 2022-10-02 014924.png b/img/Screenshot 2022-10-02 014924.png new file mode 100644 index 0000000..6d049b8 Binary files /dev/null and b/img/Screenshot 2022-10-02 014924.png differ diff --git a/img/aa.png b/img/aa.png new file mode 100644 index 0000000..97bb53c Binary files /dev/null and b/img/aa.png differ diff --git a/img/cachingchart.png b/img/cachingchart.png new file mode 100644 index 0000000..6c31d4b Binary files /dev/null and b/img/cachingchart.png differ diff --git a/img/coolmirror.png b/img/coolmirror.png new file mode 100644 index 0000000..0ea3838 Binary files /dev/null and b/img/coolmirror.png differ diff --git a/img/cornell.2022-10-10_04-17-01z.10000samp.png b/img/cornell.2022-10-10_04-17-01z.10000samp.png new file mode 100644 index 0000000..0ccd8a3 Binary files /dev/null and b/img/cornell.2022-10-10_04-17-01z.10000samp.png differ diff --git a/img/cornell.2022-10-10_04-28-41z.10000samp.png b/img/cornell.2022-10-10_04-28-41z.10000samp.png new file mode 100644 index 0000000..9fb826e Binary files /dev/null and b/img/cornell.2022-10-10_04-28-41z.10000samp.png differ diff --git a/img/cornell.2022-10-10_04-40-08z.10000samp.png b/img/cornell.2022-10-10_04-40-08z.10000samp.png new file mode 100644 index 0000000..cb6bdd8 Binary files /dev/null and b/img/cornell.2022-10-10_04-40-08z.10000samp.png differ diff --git a/img/cornell.2022-10-10_04-55-29z.10000samp.png b/img/cornell.2022-10-10_04-55-29z.10000samp.png new file mode 100644 index 0000000..7d73352 Binary files /dev/null and b/img/cornell.2022-10-10_04-55-29z.10000samp.png differ diff --git a/img/cornell.2022-10-10_05-07-44z.20000samp.png b/img/cornell.2022-10-10_05-07-44z.20000samp.png new file mode 100644 index 0000000..053fd29 Binary files /dev/null and b/img/cornell.2022-10-10_05-07-44z.20000samp.png differ diff --git a/img/cornell.2022-10-10_05-16-01z.1000samp.png b/img/cornell.2022-10-10_05-16-01z.1000samp.png new file mode 100644 index 0000000..d85857f Binary files /dev/null and b/img/cornell.2022-10-10_05-16-01z.1000samp.png differ diff --git a/img/cornell.2022-10-10_05-17-31z.10000samp.png b/img/cornell.2022-10-10_05-17-31z.10000samp.png new file mode 100644 index 0000000..78ee667 Binary files /dev/null and b/img/cornell.2022-10-10_05-17-31z.10000samp.png differ diff --git a/img/cornell.2022-10-10_05-30-40z.50000samp.png b/img/cornell.2022-10-10_05-30-40z.50000samp.png new file mode 100644 index 0000000..9fba4e5 Binary files /dev/null and b/img/cornell.2022-10-10_05-30-40z.50000samp.png differ diff --git a/img/cornell.2022-10-10_23-09-12z.5142samp.png b/img/cornell.2022-10-10_23-09-12z.5142samp.png new file mode 100644 index 0000000..cc83b12 Binary files /dev/null and b/img/cornell.2022-10-10_23-09-12z.5142samp.png differ diff --git a/img/cornell.2022-10-10_23-39-57z.1327samp.png b/img/cornell.2022-10-10_23-39-57z.1327samp.png new file mode 100644 index 0000000..c2f4ba3 Binary files /dev/null and b/img/cornell.2022-10-10_23-39-57z.1327samp.png differ diff --git a/img/cornell.2022-10-11_00-50-38z.5597samp.png b/img/cornell.2022-10-11_00-50-38z.5597samp.png new file mode 100644 index 0000000..7f32a89 Binary files /dev/null and b/img/cornell.2022-10-11_00-50-38z.5597samp.png differ diff --git a/img/cornell.2022-10-11_00-50-38z.5598samp.png b/img/cornell.2022-10-11_00-50-38z.5598samp.png new file mode 100644 index 0000000..fded08a Binary files /dev/null and b/img/cornell.2022-10-11_00-50-38z.5598samp.png differ diff --git a/img/cornell.2022-10-11_01-07-49z.5009samp.png b/img/cornell.2022-10-11_01-07-49z.5009samp.png new file mode 100644 index 0000000..c73410a Binary files /dev/null and b/img/cornell.2022-10-11_01-07-49z.5009samp.png differ diff --git a/img/cornell.2022-10-11_01-23-17z.5082samp.png b/img/cornell.2022-10-11_01-23-17z.5082samp.png new file mode 100644 index 0000000..2ffa88b Binary files /dev/null and b/img/cornell.2022-10-11_01-23-17z.5082samp.png differ diff --git a/img/cornell.2022-10-11_02-20-06z.5170samp.png b/img/cornell.2022-10-11_02-20-06z.5170samp.png new file mode 100644 index 0000000..caa87d5 Binary files /dev/null and b/img/cornell.2022-10-11_02-20-06z.5170samp.png differ diff --git a/img/cornell.2022-10-11_02-20-06z.5201samp.png b/img/cornell.2022-10-11_02-20-06z.5201samp.png new file mode 100644 index 0000000..23aa0a9 Binary files /dev/null and b/img/cornell.2022-10-11_02-20-06z.5201samp.png differ diff --git a/img/cornell.2022-10-11_02-43-13z.7754samp.png b/img/cornell.2022-10-11_02-43-13z.7754samp.png new file mode 100644 index 0000000..cea59f8 Binary files /dev/null and b/img/cornell.2022-10-11_02-43-13z.7754samp.png differ diff --git a/img/cornell.2022-10-11_02-43-13z.7759samp.png b/img/cornell.2022-10-11_02-43-13z.7759samp.png new file mode 100644 index 0000000..47ca3c8 Binary files /dev/null and b/img/cornell.2022-10-11_02-43-13z.7759samp.png differ diff --git a/img/cornell.2022-10-11_03-01-03z.11378samp.png b/img/cornell.2022-10-11_03-01-03z.11378samp.png new file mode 100644 index 0000000..2749e38 Binary files /dev/null and b/img/cornell.2022-10-11_03-01-03z.11378samp.png differ diff --git a/img/cornell.2022-10-11_03-01-03z.11379samp.png b/img/cornell.2022-10-11_03-01-03z.11379samp.png new file mode 100644 index 0000000..3117f57 Binary files /dev/null and b/img/cornell.2022-10-11_03-01-03z.11379samp.png differ diff --git a/img/cornell.2022-10-11_03-38-02z.1000samp.png b/img/cornell.2022-10-11_03-38-02z.1000samp.png new file mode 100644 index 0000000..b22ac57 Binary files /dev/null and b/img/cornell.2022-10-11_03-38-02z.1000samp.png differ diff --git a/img/cornell.2022-10-11_03-38-02z.997samp.png b/img/cornell.2022-10-11_03-38-02z.997samp.png new file mode 100644 index 0000000..fbb86f2 Binary files /dev/null and b/img/cornell.2022-10-11_03-38-02z.997samp.png differ diff --git a/img/cornell.2022-10-11_03-38-02z.999samp.png b/img/cornell.2022-10-11_03-38-02z.999samp.png new file mode 100644 index 0000000..3b912c6 Binary files /dev/null and b/img/cornell.2022-10-11_03-38-02z.999samp.png differ diff --git a/img/cornell.2022-10-11_03-40-14z.1000samp.png b/img/cornell.2022-10-11_03-40-14z.1000samp.png new file mode 100644 index 0000000..a6acc1a Binary files /dev/null and b/img/cornell.2022-10-11_03-40-14z.1000samp.png differ diff --git a/img/cornell.2022-10-11_03-53-19z.1000samp.png b/img/cornell.2022-10-11_03-53-19z.1000samp.png new file mode 100644 index 0000000..fce39af Binary files /dev/null and b/img/cornell.2022-10-11_03-53-19z.1000samp.png differ diff --git a/img/cornell.2022-10-11_03-54-58z.1000samp.png b/img/cornell.2022-10-11_03-54-58z.1000samp.png new file mode 100644 index 0000000..1c3b614 Binary files /dev/null and b/img/cornell.2022-10-11_03-54-58z.1000samp.png differ diff --git a/img/cornell.2022-10-11_03-58-36z.1000samp.png b/img/cornell.2022-10-11_03-58-36z.1000samp.png new file mode 100644 index 0000000..7862dd5 Binary files /dev/null and b/img/cornell.2022-10-11_03-58-36z.1000samp.png differ diff --git a/img/materialsortchart.png b/img/materialsortchart.png new file mode 100644 index 0000000..0698a67 Binary files /dev/null and b/img/materialsortchart.png differ diff --git a/img/materialsorttest.png b/img/materialsorttest.png new file mode 100644 index 0000000..495d9a4 Binary files /dev/null and b/img/materialsorttest.png differ diff --git a/img/noaa.png b/img/noaa.png new file mode 100644 index 0000000..eebf9e2 Binary files /dev/null and b/img/noaa.png differ diff --git a/img/thinLens8_aa.png b/img/thinLens8_aa.png new file mode 100644 index 0000000..d95cda7 Binary files /dev/null and b/img/thinLens8_aa.png differ diff --git a/img/thinlens2.png b/img/thinlens2.png new file mode 100644 index 0000000..add5ad4 Binary files /dev/null and b/img/thinlens2.png differ diff --git a/img/thinlens20.png b/img/thinlens20.png new file mode 100644 index 0000000..fae04fb Binary files /dev/null and b/img/thinlens20.png differ diff --git a/img/thinlens8.png b/img/thinlens8.png new file mode 100644 index 0000000..601985d Binary files /dev/null and b/img/thinlens8.png differ diff --git a/img/thinlens8comp.png b/img/thinlens8comp.png new file mode 100644 index 0000000..5cd241d Binary files /dev/null and b/img/thinlens8comp.png differ diff --git a/scenes/aatest.txt b/scenes/aatest.txt new file mode 100644 index 0000000..8c204c6 --- /dev/null +++ b/scenes/aatest.txt @@ -0,0 +1,166 @@ +// Red Light +MATERIAL 0 +RGB 1 1 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 10 + +// Blue Light +MATERIAL 1 +RGB 0.3 0.3 0.7 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 10 + +// Diffuse white +MATERIAL 2 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Diffuse red +MATERIAL 3 +RGB .85 .35 .35 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Diffuse green +MATERIAL 4 +RGB .35 .25 .85 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Specular white +MATERIAL 5 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB .98 .98 .98 +REFL 1 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Refractive Material +MATERIAL 6 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB .98 .98 .98 +REFL 0 +REFR 1 +REFRIOR 0 +EMITTANCE 0 + +// White Light +MATERIAL 7 +RGB 1 1 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 1 + + +// Camera +CAMERA +RES 1080 1080 +FOVY 45 +ITERATIONS 1000 +DEPTH 32 +LENSR 0 +FOCALD 7.5 +FILE cornell +EYE 0.0 7 10.5 +LOOKAT 0 7 0 +UP 0 1 0 + + +// Ceiling light +OBJECT 0 +cube +material 0 +TRANS 0 25 2 +ROTAT 0 0 0 +SCALE 4 0.1 4 +VELOCITY 0 0 0 + +// Floor +OBJECT 1 +cube +material 2 +TRANS 0 0 0 +ROTAT 0 0 0 +SCALE 100 .01 100 +VELOCITY 0 0 0 + +// Back wall +OBJECT 2 +cube +material 3 +TRANS -5 3.75 -5 +ROTAT 0 90 0 +SCALE .01 7.5 10 +VELOCITY 0 0 0 + +// Back wall +OBJECT 3 +cube +material 4 +TRANS 5 3.75 -5 +ROTAT 0 90 0 +SCALE .01 7.5 10 +VELOCITY 0 0 0 + +// Back wall +OBJECT 4 +cube +material 4 +TRANS -5 11.25 -5 +ROTAT 0 90 0 +SCALE .01 7.5 10 +VELOCITY 0 0 0 + +// Back wall +OBJECT 5 +cube +material 3 +TRANS 5 11.25 -5 +ROTAT 0 90 0 +SCALE .01 7.5 10 +VELOCITY 0 0 0 + +// Back wall +OBJECT 6 +cube +material 3 +TRANS -5 18.75 -5 +ROTAT 0 90 0 +SCALE .01 7.5 10 +VELOCITY 0 0 0 + +// Back wall +OBJECT 7 +cube +material 4 +TRANS 5 18.75 -5 +ROTAT 0 90 0 +SCALE .01 7.5 10 +VELOCITY 0 0 0 diff --git a/scenes/cornell.txt b/scenes/cornell.txt index 83ff820..3a86e11 100644 --- a/scenes/cornell.txt +++ b/scenes/cornell.txt @@ -48,12 +48,24 @@ REFR 0 REFRIOR 0 EMITTANCE 0 +// Refractive Material +MATERIAL 5 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB .98 .98 .98 +REFL 0 +REFR 1 +REFRIOR 0 +EMITTANCE 0 + // Camera CAMERA -RES 800 800 +RES 512 512 FOVY 45 ITERATIONS 5000 DEPTH 8 +LENSR 0.2 +FOCALD 8 FILE cornell EYE 0.0 5 10.5 LOOKAT 0 5 0 @@ -67,14 +79,16 @@ material 0 TRANS 0 10 0 ROTAT 0 0 0 SCALE 3 .3 3 +VELOCITY 0 0 0 // Floor OBJECT 1 cube -material 1 +material 4 TRANS 0 0 0 ROTAT 0 0 0 SCALE 10 .01 10 +VELOCITY 0 0 0 // Ceiling OBJECT 2 @@ -83,6 +97,7 @@ material 1 TRANS 0 10 0 ROTAT 0 0 90 SCALE .01 10 10 +VELOCITY 0 0 0 // Back wall OBJECT 3 @@ -91,27 +106,31 @@ material 1 TRANS 0 5 -5 ROTAT 0 90 0 SCALE .01 10 10 +VELOCITY 0 0 0 // Left wall OBJECT 4 cube -material 2 +material 3 TRANS -5 5 0 ROTAT 0 0 0 SCALE .01 10 10 +VELOCITY 0 0 0 // Right wall OBJECT 5 cube -material 3 +material 2 TRANS 5 5 0 ROTAT 0 0 0 SCALE .01 10 10 +VELOCITY 0 0 0 // Sphere OBJECT 6 sphere -material 4 -TRANS -1 4 -1 +material 5 +TRANS 0 5 0 ROTAT 0 0 0 SCALE 3 3 3 +VELOCITY 0 0 0 diff --git a/scenes/defaultsettings.txt b/scenes/defaultsettings.txt new file mode 100644 index 0000000..ffc47a9 --- /dev/null +++ b/scenes/defaultsettings.txt @@ -0,0 +1,8 @@ +// Definitions for features +#define DEBUG 0 +#define BOUNCES 4 +#define MATERIALSORT 1 +#define CACHING 0 +#define ANTIALIASING 1 +#define THINLENS 1 +#define MOTIONBLUR 1 \ No newline at end of file diff --git a/scenes/refraction.txt b/scenes/refraction.txt new file mode 100644 index 0000000..56d1e80 --- /dev/null +++ b/scenes/refraction.txt @@ -0,0 +1,220 @@ +// Red Light +MATERIAL 0 +RGB 1 0.3 0.3 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 10 + +// Blue Light +MATERIAL 1 +RGB 0.3 0.3 0.7 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 10 + +// Diffuse white +MATERIAL 2 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Diffuse red +MATERIAL 3 +RGB .85 .35 .35 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Diffuse green +MATERIAL 4 +RGB .35 .25 .85 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Specular white +MATERIAL 5 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB .98 .98 .98 +REFL 1 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Refractive Material +MATERIAL 6 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB .98 .98 .98 +REFL 0 +REFR 1 +REFRIOR 0 +EMITTANCE 0 + +// White Light +MATERIAL 7 +RGB 1 1 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 1 + + +// Camera +CAMERA +RES 1080 1080 +FOVY 45 +ITERATIONS 50000 +DEPTH 32 +LENSR 0.1 +FOCALD 7.5 +FILE cornell +EYE 0.0 5 10.5 +LOOKAT 0 5 0 +UP 0 1 0 + + +// Ceiling light +OBJECT 0 +cube +material 0 +TRANS -3 17 2 +ROTAT 0 0 0 +SCALE 4 0.1 4 +VELOCITY 0 0 0 + +// Ceiling light +OBJECT 1 +cube +material 1 +TRANS 3 17 2 +ROTAT 0 0 0 +SCALE 4 0.1 4 +VELOCITY 0 0 0 + +// Floor +OBJECT 2 +cube +material 2 +TRANS 0 0 0 +ROTAT 0 0 0 +SCALE 40 .01 100 +VELOCITY 0 0 0 + +// Left wall +OBJECT 3 +cube +material 2 +TRANS -10 5 0 +ROTAT 0 0 0 +SCALE .01 50 100 +VELOCITY 0 0 0 + +// Right wall +OBJECT 4 +cube +material 2 +TRANS 10 5 0 +ROTAT 0 0 0 +SCALE .01 50 100 +VELOCITY 0 0 0 + +// Glass Sphere +OBJECT 5 +sphere +material 6 +TRANS 0 13.5 3 +ROTAT 0 0 0 +SCALE 7 7 7 +VELOCITY 0 0 0 + +// Glass Sphere +OBJECT 6 +sphere +material 6 +TRANS 0 6.75 3 +ROTAT 0 0 0 +SCALE 5 5 5 +VELOCITY 0 0 0 + +// Glass Sphere +OBJECT 7 +sphere +material 6 +TRANS 0 2.5 3 +ROTAT 0 0 0 +SCALE 3 3 3 +VELOCITY 0 0 0 + +// Back wall +OBJECT 8 +cube +material 3 +TRANS -5 3.75 -5 +ROTAT 0 90 0 +SCALE .01 7.5 10 +VELOCITY 0 0 0 + +// Back wall +OBJECT 9 +cube +material 4 +TRANS 5 3.75 -5 +ROTAT 0 90 0 +SCALE .01 7.5 10 +VELOCITY 0 0 0 + +// Back wall +OBJECT 10 +cube +material 4 +TRANS -5 11.25 -5 +ROTAT 0 90 0 +SCALE .01 7.5 10 +VELOCITY 0 0 0 + +// Back wall +OBJECT 11 +cube +material 3 +TRANS 5 11.25 -5 +ROTAT 0 90 0 +SCALE .01 7.5 10 +VELOCITY 0 0 0 + +// Back wall +OBJECT 12 +cube +material 3 +TRANS -5 18.75 -5 +ROTAT 0 90 0 +SCALE .01 7.5 10 +VELOCITY 0 0 0 + +// Back wall +OBJECT 13 +cube +material 4 +TRANS 5 18.75 -5 +ROTAT 0 90 0 +SCALE .01 7.5 10 +VELOCITY 0 0 0 diff --git a/scenes/test.txt b/scenes/test.txt new file mode 100644 index 0000000..8d68eac --- /dev/null +++ b/scenes/test.txt @@ -0,0 +1,128 @@ +// Emissive material (light) +MATERIAL 0 +RGB 0.5 0.5 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 10 + +// 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 + +// Refractive Material +MATERIAL 5 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB .98 .98 .98 +REFL 0 +REFR 1 +REFRIOR 0 +EMITTANCE 0 + +// Emissive material (light) +MATERIAL 6 +RGB 1 0.5 0.5 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 10 + +MATERIAL 7 +RGB 1 0.5 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 10 + +// Camera +CAMERA +RES 1080 1080 +FOVY 45 +ITERATIONS 1000 +DEPTH 8 +LENSR 0 +FOCALD 3 +FILE cornell +EYE 0.0 5 10.5 +LOOKAT 0 5 0 +UP 0 1 0 + + +// Ceiling light +OBJECT 0 +sphere +material 0 +TRANS -3 17 2 +ROTAT 0 0 0 +SCALE 4 4 4 +VELOCITY 0 0 0 + +// Ceiling light +OBJECT 1 +sphere +material 0 +TRANS 3 17 2 +ROTAT 0 0 0 +SCALE 4 4 4 +VELOCITY 0 0 0 + +// Floor +OBJECT 2 +cube +material 1 +TRANS 0 0 0 +ROTAT 0 0 0 +SCALE 40 .01 40 +VELOCITY 0 0 0 + +// Sphere +OBJECT 3 +sphere +material 1 +TRANS 0 5 1.5 +ROTAT 0 0 0 +SCALE 5 5 5 +VELOCITY 0 0 0 diff --git a/scenes/thinlens.txt b/scenes/thinlens.txt new file mode 100644 index 0000000..5a3f5b0 --- /dev/null +++ b/scenes/thinlens.txt @@ -0,0 +1,192 @@ +// Emissive material (light) +MATERIAL 0 +RGB 0.5 0.5 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 10 + +// 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 + +// Refractive Material +MATERIAL 5 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB .98 .98 .98 +REFL 0 +REFR 1 +REFRIOR 0 +EMITTANCE 0 + +// Emissive material (light) +MATERIAL 6 +RGB 1 0.5 0.5 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 10 + +MATERIAL 7 +RGB 1 0.5 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 10 + +// Camera +CAMERA +RES 1080 1080 +FOVY 45 +ITERATIONS 50000 +DEPTH 32 +LENSR 0.5 +FOCALD 3 +FILE cornell +EYE 0.0 5 10.5 +LOOKAT 0 5 0 +UP 0 1 0 + + +// Ceiling light +OBJECT 0 +sphere +material 0 +TRANS -3 17 2 +ROTAT 0 0 0 +SCALE 4 4 4 +VELOCITY 0 0 0 + +// Ceiling light +OBJECT 1 +sphere +material 0 +TRANS 3 17 2 +ROTAT 0 0 0 +SCALE 4 4 4 +VELOCITY 0 0 0 + +// Floor +OBJECT 2 +cube +material 1 +TRANS 0 0 0 +ROTAT 0 0 0 +SCALE 40 .01 20 +VELOCITY 0 0 0 + +// Sphere +OBJECT 3 +sphere +material 1 +TRANS 0 5 1.5 +ROTAT 0 0 0 +SCALE 5 5 5 +VELOCITY 0 0 0 + +// Left wall +OBJECT 4 +cube +material 4 +TRANS -5 5 0 +ROTAT 45 0 0 +SCALE .01 20 20 +VELOCITY 0 0 0 + +// Right wall +OBJECT 5 +cube +material 4 +TRANS 5 5 0 +ROTAT -45 0 0 +SCALE .01 20 20 +VELOCITY 0 0 0 + +// Sphere +OBJECT 6 +sphere +material 3 +TRANS 0 15 -10 +ROTAT 0 0 0 +SCALE 5 5 5 +VELOCITY 0 0 0 + +// Sphere +OBJECT 7 +sphere +material 3 +TRANS 0 1.5 8 +ROTAT 0 0 0 +SCALE 5 5 5 +VELOCITY 0 0 0 + +// Ceiling light +OBJECT 8 +sphere +material 7 +TRANS 0 5 13.5 +ROTAT 0 0 0 +SCALE 4 4 4 +VELOCITY 0 0 0 + +// Sphere +OBJECT 9 +sphere +material 2 +TRANS 0 10 -5 +ROTAT 0 0 0 +SCALE 5 5 5 +VELOCITY 0 0 0 + +// Sphere +OBJECT 10 +sphere +material 2 +TRANS 0 2.5 5.5 +ROTAT 0 0 0 +SCALE 5 5 5 +VELOCITY 0 0 0 + diff --git a/src/interactions.h b/src/interactions.h index f969e45..4f0dd32 100644 --- a/src/interactions.h +++ b/src/interactions.h @@ -2,6 +2,8 @@ #include "intersections.h" +#define DIRECTLIGHT 1 + // CHECKITOUT /** * Computes a cosine-weighted random direction in a hemisphere. @@ -41,6 +43,29 @@ glm::vec3 calculateRandomDirectionInHemisphere( + sin(around) * over * perpendicularDirection2; } +__host__ __device__ +glm::vec3 sampleSphereLight( + glm::vec3 pos, glm::vec3 scale, thrust::default_random_engine& rng) { + thrust::uniform_real_distribution u01(0, 1); + + float phi = u01(rng) * 2.f * glm::pi(); + float theta = u01(rng) * 2.f * glm::pi(); + glm::vec3 unit_cartesian = glm::vec3(glm::cos(theta) * glm::sin(phi), + glm::sin(theta) * glm::sin(phi), + glm::cos(phi)); + return unit_cartesian * scale + pos; +} + +__host__ __device__ +glm::vec3 sampleCubeLight( + glm::vec3 pos, glm::vec3 scale, thrust::default_random_engine& rng) { + thrust::uniform_real_distribution u01(0, 1); + + glm::vec3 sample(u01(rng), u01(rng), u01(rng)); + return sample * scale + pos; +} + + /** * Scatter a ray with some probabilities according to the material properties. * For example, a diffuse surface scatters in a cosine-weighted hemisphere. @@ -66,14 +91,66 @@ glm::vec3 calculateRandomDirectionInHemisphere( * * You may need to change the parameter list for your purposes! */ + +__host__ __device__ glm::vec3 jitterRay(glm::vec3 direction, thrust::default_random_engine& rng, float radius) { + thrust::uniform_real_distribution u01(0, 1); + return glm::normalize(direction + glm::normalize(glm::vec3(u01(rng), u01(rng), u01(rng))) * radius); +} + __host__ __device__ void scatterRay( PathSegment & pathSegment, glm::vec3 intersect, 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. -} + Geom *lights, + thrust::default_random_engine &rng, + int num_lights, + int idx) { + + pathSegment.color *= m.color; + thrust::uniform_real_distribution u01(0, 1); + + if (m.hasReflective > 0.f) { + pathSegment.ray.direction = (u01(rng) < m.hasReflective) ? jitterRay(glm::reflect(pathSegment.ray.direction, normal), rng, 0.f) : calculateRandomDirectionInHemisphere(normal, rng); + } + else if (m.hasRefractive > 0.f) { + float costheta = glm::dot(normal, -pathSegment.ray.direction); + bool entering = costheta > 0; + float r_not = glm::pow((1.5 - 1) / (1.5 + 1), 2); + float r_theta = r_not + (1 - r_not) * glm::pow((1 - costheta), 5); + float eta = (entering) ? 1 / 1.5 : 1.5 / 1; + + glm::vec3 refracted = glm::refract(pathSegment.ray.direction, normal, eta); + glm::vec3 reflected = jitterRay(glm::reflect(pathSegment.ray.direction, normal), rng, 0.f); + + if (u01(rng) < r_theta || glm::abs(glm::dot(normal, refracted)) < 0.001) { + pathSegment.ray.direction = jitterRay(glm::reflect(pathSegment.ray.direction, normal), rng, 0.f); + } + else { + pathSegment.ray.direction = glm::refract(pathSegment.ray.direction, normal, eta); + } + } +#if DIRECTLIGHT == 1 + else if (pathSegment.remainingBounces == 1) { + int light_idx = glm::floor(u01(rng) * num_lights); + Geom light = lights[light_idx]; + glm::vec3 sample; + switch (light.type) { + case SPHERE: + sample = sampleSphereLight(light.translation, light.scale, rng); + break; + case CUBE: + sample = sampleCubeLight(light.translation, light.scale, rng); + break; + } + pathSegment.ray.direction = glm::normalize(sample - pathSegment.ray.origin); + } +#endif + else { + pathSegment.ray.direction = calculateRandomDirectionInHemisphere(normal, rng); + } + + pathSegment.ray.origin = intersect + pathSegment.ray.direction * 0.001f; + } + diff --git a/src/intersections.h b/src/intersections.h index b150407..e7397fe 100644 --- a/src/intersections.h +++ b/src/intersections.h @@ -142,3 +142,49 @@ __host__ __device__ float sphereIntersectionTest(Geom sphere, Ray r, return glm::length(r.origin - intersectionPoint); } + +__host__ __device__ float dummyIntersectionTest(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/main.cpp b/src/main.cpp index 96127b6..185f3be 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -107,6 +107,7 @@ void saveImage() { } void runCuda() { + if (camchanged) { iteration = 0; Camera& cam = renderState->camera; @@ -124,7 +125,9 @@ void runCuda() { cam.position = cameraPosition; cameraPosition += cam.lookAt; cam.position = cameraPosition; + camchanged = false; + } // Map OpenGL buffer object for writing from CUDA on a single GPU diff --git a/src/pathtrace.cu b/src/pathtrace.cu index fd2a464..3562847 100644 --- a/src/pathtrace.cu +++ b/src/pathtrace.cu @@ -1,9 +1,19 @@ #include #include +#include #include #include #include #include +#include +#include +#include +#include + +#include +#include +#include + #include "sceneStructs.h" #include "scene.h" @@ -14,8 +24,19 @@ #include "intersections.h" #include "interactions.h" + + #define ERRORCHECK 1 +// Definitions for features +#define DEBUG 0 +#define BOUNCES 4 +#define MATERIALSORT 0 +#define CACHING 0 +#define ANTIALIASING 1 +#define THINLENS 1 +#define MOTIONBLUR 0 + #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) { @@ -67,6 +88,34 @@ __global__ void sendImageToPBO(uchar4* pbo, glm::ivec2 resolution, } } +__device__ glm::vec2 concentricDiskSample(thrust::default_random_engine rng, float radius) +{ + thrust::uniform_real_distribution u01(0, 1); + + glm::vec2 random = glm::vec2(u01(rng), u01(rng)); + + // Map [0, 1] to [-1, 1] + glm::vec2 rOffset = random * 2.f - glm::vec2(1.f); + + // Handle degeneracy at the origin + if (rOffset == glm::vec2(0.f)) { + return glm::vec2(0.f); + } + + float r, theta; + + if (std::abs(rOffset.x) > std::abs(rOffset.y)) { + r = rOffset.x; + theta = glm::quarter_pi() * (rOffset.y / rOffset.x); + } + else { + r = rOffset.y; + theta = glm::half_pi() - glm::quarter_pi() * (rOffset.x / rOffset.y); + } + + return radius * r * glm::vec2(std::cos(theta), std::sin(theta)); +} + static Scene* hst_scene = NULL; static GuiDataContainer* guiData = NULL; static glm::vec3* dev_image = NULL; @@ -74,8 +123,10 @@ static Geom* dev_geoms = NULL; static Material* dev_materials = NULL; static PathSegment* dev_paths = NULL; static ShadeableIntersection* dev_intersections = NULL; -// TODO: static variables for device memory, any extra info you need, etc +// TODO:o static variables for device memory, any extra info you need, etc // ... +static ShadeableIntersection* dev_cache = NULL; +static Geom* dev_lights = NULL; void InitDataContainer(GuiDataContainer* imGuiData) { @@ -92,7 +143,7 @@ void pathtraceInit(Scene* scene) { cudaMemset(dev_image, 0, pixelcount * sizeof(glm::vec3)); cudaMalloc(&dev_paths, pixelcount * sizeof(PathSegment)); - + cudaMalloc(&dev_geoms, scene->geoms.size() * sizeof(Geom)); cudaMemcpy(dev_geoms, scene->geoms.data(), scene->geoms.size() * sizeof(Geom), cudaMemcpyHostToDevice); @@ -104,6 +155,12 @@ void pathtraceInit(Scene* scene) { // TODO: initialize any extra device memeory you need + cudaMalloc(&dev_cache, pixelcount * sizeof(ShadeableIntersection)); + cudaMemset(dev_cache, 0, pixelcount * sizeof(ShadeableIntersection)); + + cudaMalloc(&dev_lights, scene->lights.size() * sizeof(Geom)); + cudaMemcpy(dev_lights, scene->lights.data(), scene->lights.size() * sizeof(Geom), cudaMemcpyHostToDevice); + checkCUDAError("pathtraceInit"); } @@ -135,20 +192,53 @@ __global__ void generateRayFromCamera(Camera cam, int iter, int traceDepth, Path int index = x + (y * cam.resolution.x); PathSegment& segment = pathSegments[index]; + thrust::default_random_engine rng = makeSeededRandomEngine(iter, index, 0); + segment.ray.origin = cam.position; segment.color = glm::vec3(1.0f, 1.0f, 1.0f); // TODO: implement antialiasing by jittering the ray +#if ANTIALIASING == 1 + thrust::uniform_real_distribution u01(0, 1); + segment.ray.direction = glm::normalize(cam.view + - cam.right * cam.pixelLength.x * ((float)x + u01(rng) * cam.pixelLength.x - (float)cam.resolution.x * 0.5f) + - cam.up * cam.pixelLength.y * ((float)y + u01(rng) * cam.pixelLength.y - (float)cam.resolution.y * 0.5f) + ); +#else 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) ); +#endif +#if THINLENS == 1 + // Thin Lens Approximation + if (cam.lensRadius > 0.f) { + glm::vec2 pLens = concentricDiskSample(rng, cam.lensRadius); + glm::vec3 pFocus = segment.ray.origin + segment.ray.direction * cam.focalDistance; + segment.ray.origin += glm::vec3(pLens, 0.f); + segment.ray.direction = glm::normalize(pFocus - segment.ray.origin); + } +#endif + // antialiasing +//#if ANTIALIASING == 1 +// thrust::default_random_engine rng = makeSeededRandomEngine(iter, index, 0); +// segment.ray.direction = jitterRay(segment.ray.direction, rng, 0.01); +//#endif segment.pixelIndex = index; segment.remainingBounces = traceDepth; } } +__device__ glm::mat4 devBuildTransformationMatrix(glm::vec3 translation, glm::vec3 rotation, glm::vec3 scale) { + glm::mat4 translationMat = glm::translate(glm::mat4(), translation); + glm::mat4 rotationMat = glm::rotate(glm::mat4(), rotation.x * (float)PI / 180, glm::vec3(1, 0, 0)); + rotationMat = rotationMat * glm::rotate(glm::mat4(), rotation.y * (float)PI / 180, glm::vec3(0, 1, 0)); + rotationMat = rotationMat * glm::rotate(glm::mat4(), rotation.z * (float)PI / 180, glm::vec3(0, 0, 1)); + glm::mat4 scaleMat = glm::scale(glm::mat4(), scale); + return translationMat * rotationMat * scaleMat; +} + // TODO: // computeIntersections handles generating ray intersections ONLY. // Generating new rays is handled in your shader(s). @@ -160,6 +250,7 @@ __global__ void computeIntersections( , Geom* geoms , int geoms_size , ShadeableIntersection* intersections + , int iter ) { int path_index = blockIdx.x * blockDim.x + threadIdx.x; @@ -178,12 +269,26 @@ __global__ void computeIntersections( glm::vec3 tmp_intersect; glm::vec3 tmp_normal; + thrust::default_random_engine rng = makeSeededRandomEngine(iter, path_index, 0); + thrust::uniform_real_distribution u01(0, 1); + // naive parse through global geoms for (int i = 0; i < geoms_size; i++) { - Geom& geom = geoms[i]; +#ifdef MOTIONBLUR == 1 + Geom geom = geoms[i]; + + geom.translation += glm::vec3(u01(rng)) * geom.velocity; + + geom.transform = devBuildTransformationMatrix( + geom.translation, geom.rotation, geom.scale); + geom.inverseTransform = glm::inverse(geom.transform); + geom.invTranspose = glm::inverseTranspose(geom.transform); +#else + Geom &geom = geoms[i]; +#endif if (geom.type == CUBE) { t = boxIntersectionTest(geom, pathSegment.ray, tmp_intersect, tmp_normal, outside); @@ -193,7 +298,10 @@ __global__ void computeIntersections( t = sphereIntersectionTest(geom, pathSegment.ray, tmp_intersect, tmp_normal, outside); } // TODO: add more intersection tests here... triangle? metaball? CSG? - + else if (geom.type == DUMMY) + { + t = dummyIntersectionTest(geom, pathSegment.ray, tmp_intersect, tmp_normal, outside); + } // Compute the minimum t from the intersection tests to determine what // scene geometry object was hit first. if (t > 0.0f && t_min > t) @@ -207,7 +315,7 @@ __global__ void computeIntersections( if (hit_geom_index == -1) { - intersections[path_index].t = -1.0f; + intersections[path_index].t = -1.0f; } else { @@ -262,12 +370,14 @@ __global__ void shadeFakeMaterial( 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); } } @@ -285,6 +395,72 @@ __global__ void finalGather(int nPaths, glm::vec3* image, PathSegment* iteration } } +__global__ void shadeBSDF( + int iter + , int num_paths + , int num_lights + , ShadeableIntersection* shadeableIntersections + , PathSegment* pathSegments + , Material* materials + , Geom *lights +) +{ + 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... + // 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::uniform_real_distribution u01(0, 1); + + Material material = materials[intersection.materialId]; + glm::vec3 materialColor = material.color; + + // 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 { + scatterRay(pathSegments[idx], + getPointOnRay(pathSegments[idx].ray, intersection.t), + intersection.surfaceNormal, + material, lights, rng, num_lights, idx); + pathSegments[idx].remainingBounces--; + } + // 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].remainingBounces = 0; + } + } +} + +struct rayTerminated { + __host__ __device__ + bool operator()(const PathSegment &path) { + return path.remainingBounces > 0; + } +}; + +struct materialOrder { + __host__ __device__ + bool operator()(const ShadeableIntersection& a, const ShadeableIntersection& b){ + return a.materialId > b.materialId; + } +}; + + /** * Wrapper for the __global__ call that sets up the kernel calls and does a ton * of memory management @@ -301,36 +477,37 @@ void pathtrace(uchar4* pbo, int frame, int iter) { (cam.resolution.y + blockSize2d.y - 1) / blockSize2d.y); // 1D block for path tracing - const int blockSize1d = 128; + const int blockSize1d = 256; + const dim3 blocksPerGrid1d = ((cam.resolution.x * cam.resolution.y + blockSize1d - 1) / blockSize1d); /////////////////////////////////////////////////////////////////////////// - // 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. + //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 @@ -338,9 +515,12 @@ void pathtrace(uchar4* pbo, int frame, int iter) { checkCUDAError("generate camera ray"); int depth = 0; + PathSegment* dev_path_end = dev_paths + pixelcount; int num_paths = dev_path_end - dev_paths; + ShadeableIntersection* dev_isects = NULL; + // --- PathSegment Tracing Stage --- // Shoot ray into scene, bounce between objects, push shading chunks @@ -352,6 +532,27 @@ void pathtrace(uchar4* pbo, int frame, int iter) { // tracing dim3 numblocksPathSegmentTracing = (num_paths + blockSize1d - 1) / blockSize1d; + + + // CACHE AND LOAD FROM CACHE +#if CACHING == 1 + dev_isects = (iter != 1 && depth == 0) ? dev_cache : dev_intersections; + if (iter == 1 || depth != 0) { + computeIntersections << > > ( + depth + , num_paths + , dev_paths + , dev_geoms + , hst_scene->geoms.size() + , dev_intersections + , iter + ); + } + if (iter == 1 && depth == 0) { + cudaMemcpy(dev_cache, dev_intersections, num_paths * sizeof(ShadeableIntersection), cudaMemcpyDeviceToDevice); + } +#else + dev_isects = dev_intersections; computeIntersections << > > ( depth , num_paths @@ -359,11 +560,18 @@ void pathtrace(uchar4* pbo, int frame, int iter) { , dev_geoms , hst_scene->geoms.size() , dev_intersections + , iter ); +#endif checkCUDAError("trace one bounce"); cudaDeviceSynchronize(); depth++; +#if MATERIALSORT == 1 + thrust::sort_by_key(thrust::device, dev_intersections, dev_intersections + num_paths, dev_paths, materialOrder()); +#endif + + // TODO: // --- Shading Stage --- // Shade path segments based on intersections and generate new rays by @@ -372,16 +580,37 @@ void pathtrace(uchar4* pbo, int frame, int iter) { // 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 << > > ( +#if DEBUG == 1 + shadeBSDF << > > ( iter, num_paths, - dev_intersections, + dev_isects, dev_paths, dev_materials ); - iterationComplete = true; // TODO: should be based off stream compaction results. +#else + shadeBSDF << > > ( + iter, + num_paths, + hst_scene->lights.size(), + dev_isects, + dev_paths, + dev_materials, + dev_lights + ); +#endif + + //cudaDeviceSynchronize(); +#if DEBUG == 1 + PathSegment *new_path_end = thrust::remove_if(thrust::device, dev_paths, dev_paths + num_paths, rayTerminated()); + num_paths = new_path_end - dev_paths; + iterationComplete = (depth == BOUNCES); // TODO: should be based off stream compaction results. +#else + dev_path_end = thrust::stable_partition(thrust::device, dev_paths, dev_paths + num_paths, rayTerminated()); + num_paths = dev_path_end - dev_paths; + iterationComplete = (num_paths == 0); +#endif if (guiData != NULL) { guiData->TracedDepth = depth; @@ -390,7 +619,7 @@ void pathtrace(uchar4* pbo, int frame, int iter) { // Assemble this iteration and apply it to the image dim3 numBlocksPixels = (pixelcount + blockSize1d - 1) / blockSize1d; - finalGather << > > (num_paths, dev_image, dev_paths); + finalGather << > > (pixelcount, dev_image, dev_paths); /////////////////////////////////////////////////////////////////////////// diff --git a/src/scene.cpp b/src/scene.cpp index 3fb6239..4886f90 100644 --- a/src/scene.cpp +++ b/src/scene.cpp @@ -74,6 +74,8 @@ int Scene::loadGeom(string objectid) { newGeom.rotation = glm::vec3(atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str())); } else if (strcmp(tokens[0].c_str(), "SCALE") == 0) { newGeom.scale = glm::vec3(atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str())); + } else if (strcmp(tokens[0].c_str(), "VELOCITY") == 0) { + newGeom.velocity = glm::vec3(atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str())); } utilityCore::safeGetline(fp_in, line); @@ -85,6 +87,9 @@ int Scene::loadGeom(string objectid) { newGeom.invTranspose = glm::inverseTranspose(newGeom.transform); geoms.push_back(newGeom); + if (materials[newGeom.materialid].emittance > 0) { + lights.push_back(newGeom); + } return 1; } } @@ -96,7 +101,7 @@ int Scene::loadCamera() { float fovy; //load static properties - for (int i = 0; i < 5; i++) { + for (int i = 0; i < 7; i++) { string line; utilityCore::safeGetline(fp_in, line); vector tokens = utilityCore::tokenizeString(line); @@ -109,9 +114,13 @@ int Scene::loadCamera() { state.iterations = atoi(tokens[1].c_str()); } else if (strcmp(tokens[0].c_str(), "DEPTH") == 0) { state.traceDepth = atoi(tokens[1].c_str()); - } else if (strcmp(tokens[0].c_str(), "FILE") == 0) { + } else if (strcmp(tokens[0].c_str(), "LENSR") == 0) { + camera.lensRadius = atof(tokens[1].c_str()); + } else if (strcmp(tokens[0].c_str(), "FOCALD") == 0) { + camera.focalDistance = atof(tokens[1].c_str()); + } else if (strcmp(tokens[0].c_str(), "FILE") == 0 ) { state.imageName = tokens[1]; - } + } } string line; @@ -124,7 +133,7 @@ int Scene::loadCamera() { camera.lookAt = glm::vec3(atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str())); } 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/scene.h b/src/scene.h index f29a917..9acaf2e 100644 --- a/src/scene.h +++ b/src/scene.h @@ -22,5 +22,6 @@ class Scene { std::vector geoms; std::vector materials; + std::vector lights; RenderState state; }; diff --git a/src/sceneStructs.h b/src/sceneStructs.h index da4dbf3..37e0064 100644 --- a/src/sceneStructs.h +++ b/src/sceneStructs.h @@ -10,6 +10,8 @@ enum GeomType { SPHERE, CUBE, + + DUMMY }; struct Ray { @@ -23,6 +25,7 @@ struct Geom { glm::vec3 translation; glm::vec3 rotation; glm::vec3 scale; + glm::vec3 velocity; glm::mat4 transform; glm::mat4 inverseTransform; glm::mat4 invTranspose; @@ -49,6 +52,9 @@ struct Camera { glm::vec3 right; glm::vec2 fov; glm::vec2 pixelLength; + + float focalDistance; + float lensRadius; }; struct RenderState {