diff --git a/README.md b/README.md index 110697c..37f6c02 100644 --- a/README.md +++ b/README.md @@ -1,13 +1,115 @@ -CUDA Path Tracer -================ +####University of Pennsylvania +####CIS 565: GPU Programming and Architecture -**University of Pennsylvania, CIS 565: GPU Programming and Architecture, Project 3** +##Project 3 - CUDA Path Tracer +* Xueyin Wan +* Tested on: Windows 7, Xeon(R) E5-1630 @ 3.70GHz 32GB, GTX 1070 8192MB (Moore 103 SigLab) +* Compiled with Visual Studio 2013 and CUDA 7.5 -* (TODO) YOUR NAME HERE -* Tested on: (TODO) Windows 22, i7-2222 @ 2.22GHz 22GB, GTX 222 222MB (Moore 2222 Lab) +================================================================== +##My Path Tracer Features +###Core Features +1. A shading kernel with BSDF evaluation: + * Ideal Diffuse surfaces + * Perfectly specular-reflective (mirrored) surfaces +2. Path continuation/termination using Stream Compaction +3. toggle between sorting path/intersection continuous by material type +4. toggle between cache first bounce -### (TODO: Your README) +###Extra Coolness +1. Refraction (e.g. glass/water) [PBRT 8.2] with Frensel effects using Schlick's approximation(finally...) +*Found great reference: http://graphics.stanford.edu/courses/cs148-10-summer/docs/2006--degreve--reflection_refraction.pdf* +2. Physically-based depth-of-field +3. Motion Blur +4. More is coming! -*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. +================================================================== +###Result In Progress + +####Fresnel Refraction +Using Schlick's Approximation(http://graphics.stanford.edu/courses/cs148-10-summer/docs/2006--degreve--reflection_refraction.pdf) +![alt text](https://github.com/xueyinw/Project3-CUDA-Path-Tracer/blob/master/results/cornellTestFresnel.2016-10-09_23-18-37z.5000samp.png "Fresnel Refraction") + +####Diffuse, Perfect Specular Wall, Transmissive and Depth of Field +* `Left Wall` : Reflection = 1 +* `Left Blue Ball`: Refraction = 1 +* `Middle Purple Ball`: Refraction + Reflection = 1, No Diffuse +* `Right Pink Ball`: Only Diffuse + + +|Original | With DOF | +|------|------| +|![alt text](https://github.com/xueyinw/Project3-CUDA-Path-Tracer/blob/master/results/cornellTestDOF.2016-10-11_03-52-19z.5000samp.png "Original") | ![alt text](https://github.com/xueyinw/Project3-CUDA-Path-Tracer/blob/master/results/cornellTestDOF.2016-10-09_19-18-26z.5000samp.png "Depth Of Field and Mirror") | + +####Motion Blur +Motion Blur: The left sphere and middle cube is moving during the whole render process +![alt text](https://github.com/xueyinw/Project3-CUDA-Path-Tracer/blob/master/results/cornellTestFresnel.2016-10-10_00-02-09z.5000samp.png "Motion Blur") + +|![alt text](https://github.com/xueyinw/Project3-CUDA-Path-Tracer/blob/master/results/cornellTestFresnel.2016-10-09_18-51-14z.5000samp.png "Motion Blur") + + +####Depth of Field Comparison +* `Left Ball` : Refraction = 1 +* `Middle Cube`: Reflection = 0.4, Refraction = 0.6, Refraction + Reflection = 1, No Diffuse, +* `Right Ball`: Reflection = 0.2, Refraction = 0.2, Diffuse = 0.6 + +####Original Pic +![alt text](https://github.com/xueyinw/Project3-CUDA-Path-Tracer/blob/master/results/cornellTestFresnel.2016-10-10_00-49-35z.5000samp_original.png "Without DOF") + +|Focal Length = 10.5 | Focal Length = 9 | +|------|------| +|![alt text](https://github.com/xueyinw/Project3-CUDA-Path-Tracer/blob/master/results/cornellTestFresnel.2016-10-10_00-57-14z.5000samp_10_5.png "DOF, FOCAL LENGTH = 10.5") | ![alt text](https://github.com/xueyinw/Project3-CUDA-Path-Tracer/blob/master/results/cornellTestFresnel.2016-10-10_01-01-39z.5000samp_9.png "DOF, FOCAL LENGTH = 9") | + +####SEE WHERE I START FROM... +####BE CONFIDENT ABOUT YOUR HARD WORK!!! BE CONFIDENT ON THE WAY OF GPU!!! BE CONFIDENT WHEN FINDING A JOB!!! + +|Ideal Diffuse surfaces | Perfectly specular-reflective surfaces | +|------|------| +|![alt text](https://github.com/xueyinw/Project3-CUDA-Path-Tracer/blob/master/results/cornell.2016-10-03_01-03-54z.5000samp.png "Ideal Diffuse surfaces") | ![alt text](https://github.com/xueyinw/Project3-CUDA-Path-Tracer/blob/master/results/cornell.2016-10-03_13-08-43z.5000samp.png "Perfectly specular-reflective surfaces") | + +================================================================== +###ON or OFF? SORT_BY_MATERIAL, CACHE_FIRST_BOUNCE, STREAM_COMPATION +STATUS | SORT_BY_MATERIAL | CACHE_FIRST_BOUNCE | STREAM_COMPATION | 5000 INTERATIONS TOTAL TIME (s) +--- | --- | --- | --- | --- + | ON | OFF | OFF | 670.627 + | OFF| ON | OFF | 100.689 + | OFF| OFF | ON | 169.234 + | OFF| OFF | OFF | 117.215 +![alt text](https://github.com/xueyinw/Project3-CUDA-Path-Tracer/blob/master/results/performance_form.png "Table") + +Above results is based on my cornellTestDOFLAB.txt. +I used Siglab Machine to run all the results. + +###How to implement these Toggle? +In pathtrace.cu, using Macro definition to realize this :) +```java +#define SORT_BY_MATERIAL 0 +#define CACHE_FIRST_INTERSECTION 0 +#define STREAM_COMPACTION 0 +``` +Thanks google groups! + +For `STREAM_COMPACTION`, I used `thrust::partition()` method to calculate number of paths alive. + +For `CACHE_FIRST_INTERSECTION`, I made two helper variables: +```java +bool cache_first_intersection +static ShadeableIntersection * dev_first_intersections +``` +and use cudaMemcpy() method to copy the first batch of intersections into dev_first_intersection memory. + +For `SORT_BY_MATERIAL`, I made helper comparator `MaterialComparator()` and used `thrust::sort_by_key()` method to realize sorting by mateiral. + +###What can we learn from the table above? +* SORT_BY_MATERIAL is the slowest one.. If there're a lot of materials in this code, I think that should be better. +* CACHE_FIRST_INTERSECTION : Totally went wrong when I opened DOF. But without DOF, result is good. +* STREAM_COMPACTION : Slower than all-OFF, I think maybe the ray's amount is huge, so the method takes unnecessary time. The method itself is good, yet sorting/compaction is slow (which is a pity :(. + +###Dive into CUDA +![alt text](https://github.com/xueyinw/Project3-CUDA-Path-Tracer/blob/master/results/CUDA_SUMMARY_ONE.png "CUDA_TIMELINE") + +![alt text](https://github.com/xueyinw/Project3-CUDA-Path-Tracer/blob/master/results/CUDA_SUMMARY.PNG "CUDA_SUMMARY") +From the pictures above, we could see `ComputerIntersections` is the most time wasted(72.79%), since this function is used to calculate all the intersections, and as we noticed, the ray's amount is huge. + +The second part is `shadeFakeMaterial`, since we need to deal with each ray's color, remainging bounces information. diff --git a/results/CUDA_SUMMARY.PNG b/results/CUDA_SUMMARY.PNG new file mode 100644 index 0000000..c9c26cc Binary files /dev/null and b/results/CUDA_SUMMARY.PNG differ diff --git a/results/CUDA_SUMMARY_ONE.png b/results/CUDA_SUMMARY_ONE.png new file mode 100644 index 0000000..e32470f Binary files /dev/null and b/results/CUDA_SUMMARY_ONE.png differ diff --git a/results/cornell.2016-10-03_01-03-54z.5000samp.png b/results/cornell.2016-10-03_01-03-54z.5000samp.png new file mode 100644 index 0000000..6703dd2 Binary files /dev/null and b/results/cornell.2016-10-03_01-03-54z.5000samp.png differ diff --git a/results/cornell.2016-10-03_13-08-43z.5000samp.png b/results/cornell.2016-10-03_13-08-43z.5000samp.png new file mode 100644 index 0000000..a243a05 Binary files /dev/null and b/results/cornell.2016-10-03_13-08-43z.5000samp.png differ diff --git a/results/cornell.2016-10-09_03-39-16z.2452samp.png b/results/cornell.2016-10-09_03-39-16z.2452samp.png new file mode 100644 index 0000000..26d5298 Binary files /dev/null and b/results/cornell.2016-10-09_03-39-16z.2452samp.png differ diff --git a/results/cornell.2016-10-09_05-48-25z.2806samp.png b/results/cornell.2016-10-09_05-48-25z.2806samp.png new file mode 100644 index 0000000..2a92430 Binary files /dev/null and b/results/cornell.2016-10-09_05-48-25z.2806samp.png differ diff --git a/results/cornellTestDOF.2016-10-09_19-18-26z.5000samp.png b/results/cornellTestDOF.2016-10-09_19-18-26z.5000samp.png new file mode 100644 index 0000000..3e43d97 Binary files /dev/null and b/results/cornellTestDOF.2016-10-09_19-18-26z.5000samp.png differ diff --git a/results/cornellTestDOF.2016-10-11_03-52-19z.5000samp.png b/results/cornellTestDOF.2016-10-11_03-52-19z.5000samp.png new file mode 100644 index 0000000..c9b12c8 Binary files /dev/null and b/results/cornellTestDOF.2016-10-11_03-52-19z.5000samp.png differ diff --git a/results/cornellTestFresnel.2016-10-09_12-04-53z.5000samp.png b/results/cornellTestFresnel.2016-10-09_12-04-53z.5000samp.png new file mode 100644 index 0000000..e1e8350 Binary files /dev/null and b/results/cornellTestFresnel.2016-10-09_12-04-53z.5000samp.png differ diff --git a/results/cornellTestFresnel.2016-10-09_15-28-33z.1038samp.png b/results/cornellTestFresnel.2016-10-09_15-28-33z.1038samp.png new file mode 100644 index 0000000..5d04a12 Binary files /dev/null and b/results/cornellTestFresnel.2016-10-09_15-28-33z.1038samp.png differ diff --git a/results/cornellTestFresnel.2016-10-09_15-28-33z.1294samp.png b/results/cornellTestFresnel.2016-10-09_15-28-33z.1294samp.png new file mode 100644 index 0000000..ff5ac54 Binary files /dev/null and b/results/cornellTestFresnel.2016-10-09_15-28-33z.1294samp.png differ diff --git a/results/cornellTestFresnel.2016-10-09_15-28-33z.1426samp.png b/results/cornellTestFresnel.2016-10-09_15-28-33z.1426samp.png new file mode 100644 index 0000000..523f11d Binary files /dev/null and b/results/cornellTestFresnel.2016-10-09_15-28-33z.1426samp.png differ diff --git a/results/cornellTestFresnel.2016-10-09_15-28-33z.1820samp.png b/results/cornellTestFresnel.2016-10-09_15-28-33z.1820samp.png new file mode 100644 index 0000000..5807382 Binary files /dev/null and b/results/cornellTestFresnel.2016-10-09_15-28-33z.1820samp.png differ diff --git a/results/cornellTestFresnel.2016-10-09_15-28-33z.2006samp.png b/results/cornellTestFresnel.2016-10-09_15-28-33z.2006samp.png new file mode 100644 index 0000000..784133b Binary files /dev/null and b/results/cornellTestFresnel.2016-10-09_15-28-33z.2006samp.png differ diff --git a/results/cornellTestFresnel.2016-10-09_18-51-14z.5000samp.png b/results/cornellTestFresnel.2016-10-09_18-51-14z.5000samp.png new file mode 100644 index 0000000..3ebccfa Binary files /dev/null and b/results/cornellTestFresnel.2016-10-09_18-51-14z.5000samp.png differ diff --git a/results/cornellTestFresnel.2016-10-09_23-18-37z.5000samp.png b/results/cornellTestFresnel.2016-10-09_23-18-37z.5000samp.png new file mode 100644 index 0000000..16b54bc Binary files /dev/null and b/results/cornellTestFresnel.2016-10-09_23-18-37z.5000samp.png differ diff --git a/results/cornellTestFresnel.2016-10-10_00-02-09z.5000samp.png b/results/cornellTestFresnel.2016-10-10_00-02-09z.5000samp.png new file mode 100644 index 0000000..e1ceee1 Binary files /dev/null and b/results/cornellTestFresnel.2016-10-10_00-02-09z.5000samp.png differ diff --git a/results/cornellTestFresnel.2016-10-10_00-49-35z.5000samp_original.png b/results/cornellTestFresnel.2016-10-10_00-49-35z.5000samp_original.png new file mode 100644 index 0000000..d9b3681 Binary files /dev/null and b/results/cornellTestFresnel.2016-10-10_00-49-35z.5000samp_original.png differ diff --git a/results/cornellTestFresnel.2016-10-10_00-57-14z.5000samp_10_5.png b/results/cornellTestFresnel.2016-10-10_00-57-14z.5000samp_10_5.png new file mode 100644 index 0000000..0cb14eb Binary files /dev/null and b/results/cornellTestFresnel.2016-10-10_00-57-14z.5000samp_10_5.png differ diff --git a/results/cornellTestFresnel.2016-10-10_01-01-39z.5000samp_9.png b/results/cornellTestFresnel.2016-10-10_01-01-39z.5000samp_9.png new file mode 100644 index 0000000..7f4c87f Binary files /dev/null and b/results/cornellTestFresnel.2016-10-10_01-01-39z.5000samp_9.png differ diff --git a/results/performance_form.png b/results/performance_form.png new file mode 100644 index 0000000..d27d232 Binary files /dev/null and b/results/performance_form.png differ diff --git a/scenes/cornell.txt b/scenes/cornell.txt index 83ff820..3b0a7fd 100644 --- a/scenes/cornell.txt +++ b/scenes/cornell.txt @@ -43,9 +43,9 @@ MATERIAL 4 RGB .98 .98 .98 SPECEX 0 SPECRGB .98 .98 .98 -REFL 1 -REFR 0 -REFRIOR 0 +REFL 0 +REFR 1 +REFRIOR 2.75 EMITTANCE 0 // Camera @@ -58,7 +58,8 @@ FILE cornell EYE 0.0 5 10.5 LOOKAT 0 5 0 UP 0 1 0 - +LENSRADIUS 1 +FOCALDISTANCE 10.5 // Ceiling light OBJECT 0 @@ -112,6 +113,15 @@ SCALE .01 10 10 OBJECT 6 sphere material 4 -TRANS -1 4 -1 +TRANS -3 4 -1 ROTAT 0 0 0 SCALE 3 3 3 + +// Sphere +OBJECT 7 +sphere +material 4 +TRANS 3 4 -1 +ROTAT 0 0 0 +SCALE 4 4 4 + diff --git a/scenes/cornellTestDOF.txt b/scenes/cornellTestDOF.txt new file mode 100644 index 0000000..c489b49 --- /dev/null +++ b/scenes/cornellTestDOF.txt @@ -0,0 +1,154 @@ +// 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 + +MATERIAL 4 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB .98 .98 .98 +REFL 0 +REFR 1 +REFRIOR 2.75 +EMITTANCE 0 + +// Diffuse +MATERIAL 5 +RGB 0.8 0.5 0.2 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +MATERIAL 6 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB .98 .98 .98 +REFL 0.2 +REFR 0.2 +REFRIOR 2.5 +EMITTANCE 0 + +// Camera +CAMERA +RES 800 800 +FOVY 45 +ITERATIONS 5000 +DEPTH 8 +FILE cornellTestDOF +EYE 0.0 5 10.5 +LOOKAT 0 5 0 +UP 0 1 0 +LENSRADIUS 1 +FOCALDISTANCE 9 + + +// 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 6 +TRANS -3 4 -1 +ROTAT 0 0 0 +SCALE 3 3 3 + +// Sphere +OBJECT 7 +sphere +material 6 +TRANS 3 4 -1 +ROTAT 0 0 0 +SCALE 4 4 4 + +// sphere +OBJECT 8 +sphere +material 4 +TRANS 1 5 0 +ROTAT 0 0 0 +SCALE 3 3 3 + diff --git a/scenes/cornellTestDOFLAB.txt b/scenes/cornellTestDOFLAB.txt new file mode 100644 index 0000000..86208ac --- /dev/null +++ b/scenes/cornellTestDOFLAB.txt @@ -0,0 +1,175 @@ +// Emissive material (light) +MATERIAL 0 +RGB 1 1 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 15 + +// 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 .98 .98 .98 +SPECEX 1 +SPECRGB 1 1 1 +REFL 1 +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 + +// TransMissive +MATERIAL 4 +RGB 0.3 0.6 0.9 +SPECEX 0 +SPECRGB .98 .98 .98 +REFL 0 +REFR 1 +REFRIOR 2.75 +EMITTANCE 0 + +// Diffuse +MATERIAL 5 +RGB 0.8 0.5 0.5 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Reflective + Refractive +MATERIAL 6 +RGB 1 0.3 0.8 +SPECEX 0 +SPECRGB .98 .98 .98 +REFL 0.8 +REFR 0.2 +REFRIOR 1.8 +EMITTANCE 0 + + +MATERIAL 7 +RGB .98 .98 .98 +SPECEX 20 +SPECRGB 1 1 1 +REFL 1 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +MATERIAL 8 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 1 +REFRIOR 2.5 +EMITTANCE 0 + +// Camera +CAMERA +RES 800 800 +FOVY 45 +ITERATIONS 5000 +DEPTH 8 +FILE cornellTestDOF +EYE 0.0 5 10.5 +LOOKAT 0 5 0 +UP 0 1 0 +LENSRADIUS 1 +FOCALDISTANCE 11.5 + + +// Ceiling light +OBJECT 0 +cube +material 0 +TRANS 0 10 0 +ROTAT 0 0 0 +SCALE 3 .3 3 + +// Floor +OBJECT 1 +cube +material 1 +TRANS 0 0 0 +ROTAT 0 0 0 +SCALE 10 .01 10 + +// Ceiling +OBJECT 2 +cube +material 1 +TRANS 0 10 0 +ROTAT 0 0 90 +SCALE .01 10 10 + +// Back wall +OBJECT 3 +cube +material 1 +TRANS 0 5 -5 +ROTAT 0 90 0 +SCALE .01 10 10 + +// Left wall +OBJECT 4 +cube +material 2 +TRANS -5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +// Right wall +OBJECT 5 +cube +material 3 +TRANS 5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +// Sphere +OBJECT 6 +sphere +material 4 +TRANS -3 4 -1 +ROTAT 0 0 0 +SCALE 3 3 3 + +// Sphere +OBJECT 7 +sphere +material 5 +TRANS 4 4 -1 +ROTAT 0 0 0 +SCALE 2 2 2 + +// sphere +OBJECT 8 +sphere +material 6 +TRANS 1 5 -2 +ROTAT 0 0 0 +SCALE 3 3 3 + diff --git a/scenes/cornellTestFresnel.txt b/scenes/cornellTestFresnel.txt new file mode 100644 index 0000000..08fd0a5 --- /dev/null +++ b/scenes/cornellTestFresnel.txt @@ -0,0 +1,164 @@ +// 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 + +MATERIAL 4 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB .98 .98 .98 +REFL 0.1 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +MATERIAL 5 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB .25 .75 .75 +REFL 1 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +MATERIAL 6 +RGB .88 .58 .58 +SPECEX 0 +SPECRGB .98 .98 .98 +REFL 0.8 +REFR 0.2 +REFRIOR 2.5 +EMITTANCE 0 + +MATERIAL 7 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB .98 .98 .98 +REFL 0 +REFR 1 +REFRIOR 2.5 +EMITTANCE 0 + +// Camera +CAMERA +RES 800 800 +FOVY 45 +ITERATIONS 5000 +DEPTH 8 +FILE cornellTestFresnel +EYE 0.0 5 10.5 +LOOKAT 0 5 0 +UP 0 1 0 +LENSRADIUS 1 +FOCALDISTANCE 11 + +// 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 -3 4 -1 +ROTAT 0 0 0 +SCALE 3 3 3 +MOTION 0 3 0 + +// cube +OBJECT 7 +cube +material 6 +TRANS 2 2 1 +ROTAT 0 60 0 +SCALE 2 2 2 +MOTION 2 0 0 + +// Sphere +OBJECT 8 +sphere +material 7 +TRANS 2.5 5.5 1 +ROTAT 0 0 0 +SCALE 3 3 3 +MOTION 0 0 0 \ No newline at end of file diff --git a/scenes/cornellTestFresnelLAB.txt b/scenes/cornellTestFresnelLAB.txt new file mode 100644 index 0000000..6780187 --- /dev/null +++ b/scenes/cornellTestFresnelLAB.txt @@ -0,0 +1,164 @@ +// Emissive material (light) +MATERIAL 0 +RGB 1 1 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 15 + + +// 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 + +MATERIAL 4 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB .98 .98 .98 +REFL 0.1 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +MATERIAL 5 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB .25 .25 .8 +REFL 0.3 +REFR 0.2 +REFRIOR 2.5 +EMITTANCE 0 + +MATERIAL 6 +RGB .88 .58 .58 +SPECEX 0 +SPECRGB .98 .98 .98 +REFL 0.8 +REFR 0.2 +REFRIOR 2.5 +EMITTANCE 0 + +MATERIAL 7 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 1 +REFRIOR 2.5 +EMITTANCE 0 + +// Camera +CAMERA +RES 800 800 +FOVY 45 +ITERATIONS 5000 +DEPTH 8 +FILE cornellTestFresnel +EYE 0.0 5 10.5 +LOOKAT 0 5 0 +UP 0 1 0 +LENSRADIUS 1 +FOCALDISTANCE 10 + +// 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 -3 4 -1 +ROTAT 0 30 20 +SCALE 3 3 3 +MOTION 0 1.5 0 + +// cube +OBJECT 7 +cube +material 6 +TRANS -1 2 3 +ROTAT 0 60 30 +SCALE 2 2 2 +MOTION -1 0 0 + +// Sphere +OBJECT 8 +sphere +material 7 +TRANS 2.5 5.5 2 +ROTAT 0 0 0 +SCALE 2 2 2 +MOTION 0 0 0 \ No newline at end of file diff --git a/scenes/cornellTestMirror.txt b/scenes/cornellTestMirror.txt new file mode 100644 index 0000000..0553390 --- /dev/null +++ b/scenes/cornellTestMirror.txt @@ -0,0 +1,156 @@ +// Emissive material (light) +MATERIAL 0 +RGB 1 1 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 15 + +// 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 .98 .98 .98 +SPECEX 1 +SPECRGB 1 1 1 +REFL 1 +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 + +// TransMissive +MATERIAL 4 +RGB 0.3 0.6 0.9 +SPECEX 0 +SPECRGB .98 .98 .98 +REFL 0 +REFR 1 +REFRIOR 2.75 +EMITTANCE 0 + +// Diffuse +MATERIAL 5 +RGB 0.8 0.5 0.5 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Reflective + Refractive +MATERIAL 6 +RGB 0.6 0.3 0.7 +SPECEX 0 +SPECRGB .98 .98 .98 +REFL 0.2 +REFR 0.8 +REFRIOR 2.5 +EMITTANCE 0 + +// Camera +CAMERA +RES 800 800 +FOVY 45 +ITERATIONS 5000 +DEPTH 8 +FILE cornellTestDOF +EYE 0.0 5 10.5 +LOOKAT 0 5 0 +UP 0 1 0 +LENSRADIUS 1 +FOCALDISTANCE 11 + + +// Ceiling light +OBJECT 0 +cube +material 0 +TRANS 0 10 0 +ROTAT 0 0 0 +SCALE 3 .3 3 + +// Floor +OBJECT 1 +cube +material 1 +TRANS 0 0 0 +ROTAT 0 0 0 +SCALE 10 .01 10 + +// Ceiling +OBJECT 2 +cube +material 1 +TRANS 0 10 0 +ROTAT 0 0 90 +SCALE .01 10 10 + +// Back wall +OBJECT 3 +cube +material 1 +TRANS 0 5 -5 +ROTAT 0 90 0 +SCALE .01 10 10 + +// Left wall +OBJECT 4 +cube +material 2 +TRANS -5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +// Right wall +OBJECT 5 +cube +material 3 +TRANS 5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +// Sphere +OBJECT 6 +sphere +material 4 +TRANS -3 4 -1 +ROTAT 0 0 0 +SCALE 3 3 3 + +// Sphere +OBJECT 7 +sphere +material 5 +TRANS 3 4 1 +ROTAT 0 0 0 +SCALE 2 2 2 + +// sphere +OBJECT 8 +sphere +material 6 +TRANS 1 5 -3 +ROTAT 0 0 0 +SCALE 3 3 3 + diff --git a/scenes/sphere.txt b/scenes/sphere.txt index a74b545..c475b81 100644 --- a/scenes/sphere.txt +++ b/scenes/sphere.txt @@ -18,6 +18,8 @@ FILE sphere EYE 0.0 5 10.5 LOOKAT 0 5 0 UP 0 1 0 +LENSRADIUS 1 +FOCALDISTANCE 10.5 // Sphere OBJECT 0 diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index a1cb3fb..eb47749 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_30 ) diff --git a/src/image.h b/src/image.h index ae1b768..9c4f182 100644 --- a/src/image.h +++ b/src/image.h @@ -2,7 +2,7 @@ #include -using namespace std; +//using namespace std; class image { private: diff --git a/src/interactions.h b/src/interactions.h index 5ce3628..61c5c1f 100644 --- a/src/interactions.h +++ b/src/interactions.h @@ -1,7 +1,7 @@ -#pragma once +#pragma once #include "intersections.h" - + // CHECKITOUT /** * Computes a cosine-weighted random direction in a hemisphere. @@ -50,7 +50,7 @@ glm::vec3 calculateRandomDirectionInHemisphere( * * The visual effect you want is to straight-up add the diffuse and specular * components. You can do this in a few ways. This logic also applies to - * combining other types of materias (such as refractive). + * combining other types of materials (such as refractive). * * - Always take an even (50/50) split between a each effect (a diffuse bounce * and a specular bounce), but divide the resulting color of either branch @@ -68,12 +68,70 @@ glm::vec3 calculateRandomDirectionInHemisphere( */ __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. +PathSegment & pathSegment, +glm::vec3 intersect, +glm::vec3 normal, +const Material &m, +thrust::default_random_engine &rng) { + thrust::uniform_real_distribution u01(0, 1); + float prob = u01(rng); + //my reference:http://graphics.stanford.edu/courses/cs148-10-summer/docs/2006--degreve--reflection_refraction.pdf + if (prob > 1 - m.hasRefractive && prob < 1) { //transmissive dominate + float indexRatio; + float theta = (180.0f / PI) * acos(glm::dot(pathSegment.ray.direction, normal) / (glm::length(pathSegment.ray.direction) * glm::length(normal))); + if (theta >= 90.0f) { + indexRatio = 1.f / m.indexOfRefraction; + } else { + indexRatio = m.indexOfRefraction; + } + float R0 = (1 - indexRatio) / (1 + indexRatio) * (1 - indexRatio) / (1 + indexRatio); + // Schlick’s approximation of the Fresnel equation + float RSchlick = R0 + (1.0f - R0) * glm::pow(1.0f - glm::abs(glm::dot(normal, pathSegment.ray.direction)), 5); + if (RSchlick < prob) { // refraction + pathSegment.ray.direction = glm::normalize(glm::refract(pathSegment.ray.direction, normal, indexRatio)); + } else { // reflection + pathSegment.ray.direction = glm::normalize(glm::reflect(pathSegment.ray.direction, normal)); + } + pathSegment.ray.origin = intersect + 1e-3f * (glm::normalize(pathSegment.ray.direction)); + pathSegment.color *= m.color * m.specular.color; + + } else if (m.hasReflective == 1) { //directly goes to perfect specular + pathSegment.ray.direction = pathSegment.ray.direction - 2.0f * normal * (glm::dot(pathSegment.ray.direction, normal)); + pathSegment.ray.origin = intersect + 1e-3f * (glm::normalize(pathSegment.ray.direction)); + pathSegment.color *= m.color * m.specular.color; + + } else if (prob > (1 - m.hasRefractive - m.hasReflective) && prob < (1 - m.hasRefractive)) { //reflection dominate & diffuse combined + if (0.5f * m.hasReflective < u01(rng)) { //50% percent + pathSegment.ray.direction = pathSegment.ray.direction - 2.0f * normal * (glm::dot(pathSegment.ray.direction, normal)); + pathSegment.ray.origin = intersect + 1e-3f * (glm::normalize(pathSegment.ray.direction)); + pathSegment.color *= m.color * m.specular.color; + } else { // diffuse + pathSegment.ray.direction = calculateRandomDirectionInHemisphere(normal, rng); + pathSegment.ray.origin = intersect + 1e-3f * (glm::normalize(pathSegment.ray.direction)); + pathSegment.color *= m.color; + } + + } else { + pathSegment.ray.direction = calculateRandomDirectionInHemisphere(normal, rng); + pathSegment.ray.origin = intersect + 1e-3f * (glm::normalize(pathSegment.ray.direction)); + pathSegment.color *= m.color; + } + pathSegment.remainingBounces -= 1; } +//**********my original non-frensel refraction************// + //glm::vec3 refractionDirection; + //float theta; + ////get the angel between ray.direction and the normal. + ////refer to pbrt & adam's slides + //theta = (180.0f / PI) * acos(glm::dot(pathSegment.ray.direction, normal) / (glm::length(pathSegment.ray.direction) * glm::length(normal))); + ////For the incident vector I and surface normal N, and the ratio of indices of refraction eta, return the refraction vector. + //if (theta >= 90.0f) { + // refractionDirection = glm::refract(pathSegment.ray.direction, normal, 1.0f / m.indexOfRefraction); + //} else{ + // refractionDirection = glm::refract(pathSegment.ray.direction, -normal, m.indexOfRefraction); + //} + //pathSegment.ray.direction = refractionDirection; + //pathSegment.ray.origin = intersect + glm::normalize(pathSegment.ray.direction) * 1e-3f; + //pathSegment.color *= m.color * m.specular.color; + + diff --git a/src/main.cpp b/src/main.cpp index fe8e85e..e9df802 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -1,7 +1,7 @@ #include "main.h" #include "preview.h" #include - +#include static std::string startTimeString; // For camera controls @@ -26,180 +26,194 @@ int iteration; int width; int height; +//test time +std::chrono::high_resolution_clock::time_point beginTime; +std::chrono::high_resolution_clock::time_point endTime; + //------------------------------- //-------------MAIN-------------- //------------------------------- int main(int argc, char** argv) { - startTimeString = currentTimeString(); + startTimeString = currentTimeString(); - if (argc < 2) { - printf("Usage: %s SCENEFILE.txt\n", argv[0]); - return 1; - } + if (argc < 2) { + printf("Usage: %s SCENEFILE.txt\n", argv[0]); + return 1; + } - const char *sceneFile = argv[1]; + const char *sceneFile = argv[1]; - // Load scene file - scene = new Scene(sceneFile); + // Load scene file + scene = new Scene(sceneFile); - // Set up camera stuff from loaded path tracer settings - iteration = 0; - renderState = &scene->state; - Camera &cam = renderState->camera; - width = cam.resolution.x; - height = cam.resolution.y; + // Set up camera stuff from loaded path tracer settings + iteration = 0; + renderState = &scene->state; + Camera &cam = renderState->camera; + width = cam.resolution.x; + height = cam.resolution.y; - glm::vec3 view = cam.view; - glm::vec3 up = cam.up; - glm::vec3 right = glm::cross(view, up); - up = glm::cross(right, view); + glm::vec3 view = cam.view; + glm::vec3 up = cam.up; + glm::vec3 right = glm::cross(view, up); + up = glm::cross(right, view); - cameraPosition = cam.position; + cameraPosition = cam.position; - // compute phi (horizontal) and theta (vertical) relative 3D axis - // so, (0 0 1) is forward, (0 1 0) is up - glm::vec3 viewXZ = glm::vec3(view.x, 0.0f, view.z); - glm::vec3 viewZY = glm::vec3(0.0f, view.y, view.z); - phi = glm::acos(glm::dot(glm::normalize(viewXZ), glm::vec3(0, 0, -1))); - theta = glm::acos(glm::dot(glm::normalize(viewZY), glm::vec3(0, 1, 0))); - ogLookAt = cam.lookAt; - zoom = glm::length(cam.position - ogLookAt); + // compute phi (horizontal) and theta (vertical) relative 3D axis + // so, (0 0 1) is forward, (0 1 0) is up + glm::vec3 viewXZ = glm::vec3(view.x, 0.0f, view.z); + glm::vec3 viewZY = glm::vec3(0.0f, view.y, view.z); + phi = glm::acos(glm::dot(glm::normalize(viewXZ), glm::vec3(0, 0, -1))); + theta = glm::acos(glm::dot(glm::normalize(viewZY), glm::vec3(0, 1, 0))); + ogLookAt = cam.lookAt; + zoom = glm::length(cam.position - ogLookAt); - // Initialize CUDA and GL components - init(); + // Initialize CUDA and GL components + init(); - // GLFW main loop - mainLoop(); + // GLFW main loop + mainLoop(); - return 0; + return 0; } void saveImage() { - float samples = iteration; - // output image file - image img(width, height); - - for (int x = 0; x < width; x++) { - for (int y = 0; y < height; y++) { - int index = x + (y * width); - glm::vec3 pix = renderState->image[index]; - img.setPixel(width - 1 - x, y, glm::vec3(pix) / samples); - } - } - - std::string filename = renderState->imageName; - std::ostringstream ss; - ss << filename << "." << startTimeString << "." << samples << "samp"; - filename = ss.str(); - - // CHECKITOUT - img.savePNG(filename); - //img.saveHDR(filename); // Save a Radiance HDR file + float samples = iteration; + // output image file + image img(width, height); + + for (int x = 0; x < width; x++) { + for (int y = 0; y < height; y++) { + int index = x + (y * width); + glm::vec3 pix = renderState->image[index]; + img.setPixel(width - 1 - x, y, glm::vec3(pix) / samples); + } + } + + std::string filename = renderState->imageName; + std::ostringstream ss; + ss << filename << "." << startTimeString << "." << samples << "samp"; + filename = ss.str(); + + // CHECKITOUT + img.savePNG(filename); + //img.saveHDR(filename); // Save a Radiance HDR file } void runCuda() { - if (camchanged) { - iteration = 0; - Camera &cam = renderState->camera; - cameraPosition.x = zoom * sin(phi) * sin(theta); - cameraPosition.y = zoom * cos(theta); - cameraPosition.z = zoom * cos(phi) * sin(theta); - - cam.view = -glm::normalize(cameraPosition); - glm::vec3 v = cam.view; - glm::vec3 u = glm::vec3(0, 1, 0);//glm::normalize(cam.up); - glm::vec3 r = glm::cross(v, u); - cam.up = glm::cross(r, v); - cam.right = r; - - cam.position = cameraPosition; - cameraPosition += cam.lookAt; - cam.position = cameraPosition; - camchanged = false; - } - - // Map OpenGL buffer object for writing from CUDA on a single GPU - // No data is moved (Win & Linux). When mapped to CUDA, OpenGL should not use this buffer - - if (iteration == 0) { - pathtraceFree(); - pathtraceInit(scene); - } - - if (iteration < renderState->iterations) { - uchar4 *pbo_dptr = NULL; - iteration++; - cudaGLMapBufferObject((void**)&pbo_dptr, pbo); - - // execute the kernel - int frame = 0; - pathtrace(pbo_dptr, frame, iteration); - - // unmap buffer object - cudaGLUnmapBufferObject(pbo); - } else { - saveImage(); - pathtraceFree(); - cudaDeviceReset(); - exit(EXIT_SUCCESS); - } + if (camchanged) { + iteration = 0; + Camera &cam = renderState->camera; + cameraPosition.x = zoom * sin(phi) * sin(theta); + cameraPosition.y = zoom * cos(theta); + cameraPosition.z = zoom * cos(phi) * sin(theta); + + cam.view = -glm::normalize(cameraPosition); + glm::vec3 v = cam.view; + glm::vec3 u = glm::vec3(0, 1, 0);//glm::normalize(cam.up); + glm::vec3 r = glm::cross(v, u); + cam.up = glm::cross(r, v); + cam.right = r; + + cam.position = cameraPosition; + cameraPosition += cam.lookAt; + cam.position = cameraPosition; + camchanged = false; + + beginTime = std::chrono::high_resolution_clock::now(); + } + + + // Map OpenGL buffer object for writing from CUDA on a single GPU + // No data is moved (Win & Linux). When mapped to CUDA, OpenGL should not use this buffer + + if (iteration == 0) { + pathtraceFree(); + pathtraceInit(scene); + } + + if (iteration < renderState->iterations) { + uchar4 *pbo_dptr = NULL; + iteration++; + cudaGLMapBufferObject((void**)&pbo_dptr, pbo); + + // execute the kernel + int frame = 0; + pathtrace(pbo_dptr, frame, iteration); + + // unmap buffer object + cudaGLUnmapBufferObject(pbo); + } + else { + endTime = std::chrono::high_resolution_clock::now(); + std::chrono::duration duro = endTime - beginTime; + float deltaTime = duro.count(); + double oneIterationTime = deltaTime / renderState->iterations; + std::cout << " Total time " << deltaTime / 1000 << "s" << std::endl; + std::cout << " Each Interation takes " << oneIterationTime << "ms" << std::endl; + saveImage(); + pathtraceFree(); + cudaDeviceReset(); + exit(EXIT_SUCCESS); + } } void keyCallback(GLFWwindow* window, int key, int scancode, int action, int mods) { - if (action == GLFW_PRESS) { - switch (key) { - case GLFW_KEY_ESCAPE: - saveImage(); - glfwSetWindowShouldClose(window, GL_TRUE); - break; - case GLFW_KEY_S: - saveImage(); - break; - case GLFW_KEY_SPACE: - camchanged = true; - renderState = &scene->state; - Camera &cam = renderState->camera; - cam.lookAt = ogLookAt; - break; - } - } + if (action == GLFW_PRESS) { + switch (key) { + case GLFW_KEY_ESCAPE: + saveImage(); + glfwSetWindowShouldClose(window, GL_TRUE); + break; + case GLFW_KEY_S: + saveImage(); + break; + case GLFW_KEY_SPACE: + camchanged = true; + renderState = &scene->state; + Camera &cam = renderState->camera; + cam.lookAt = ogLookAt; + break; + } + } } void mouseButtonCallback(GLFWwindow* window, int button, int action, int mods) { - leftMousePressed = (button == GLFW_MOUSE_BUTTON_LEFT && action == GLFW_PRESS); - rightMousePressed = (button == GLFW_MOUSE_BUTTON_RIGHT && action == GLFW_PRESS); - middleMousePressed = (button == GLFW_MOUSE_BUTTON_MIDDLE && action == GLFW_PRESS); + leftMousePressed = (button == GLFW_MOUSE_BUTTON_LEFT && action == GLFW_PRESS); + rightMousePressed = (button == GLFW_MOUSE_BUTTON_RIGHT && action == GLFW_PRESS); + middleMousePressed = (button == GLFW_MOUSE_BUTTON_MIDDLE && action == GLFW_PRESS); } void mousePositionCallback(GLFWwindow* window, double xpos, double ypos) { - if (xpos == lastX || ypos == lastY) return; // otherwise, clicking back into window causes re-start - if (leftMousePressed) { - // compute new camera parameters - phi -= (xpos - lastX) / width; - theta -= (ypos - lastY) / height; - theta = std::fmax(0.001f, std::fmin(theta, PI)); - camchanged = true; - } - else if (rightMousePressed) { - zoom += (ypos - lastY) / height; - zoom = std::fmax(0.1f, zoom); - camchanged = true; - } - else if (middleMousePressed) { - renderState = &scene->state; - Camera &cam = renderState->camera; - glm::vec3 forward = cam.view; - forward.y = 0.0f; - forward = glm::normalize(forward); - glm::vec3 right = cam.right; - right.y = 0.0f; - right = glm::normalize(right); - - cam.lookAt -= (float) (xpos - lastX) * right * 0.01f; - cam.lookAt += (float) (ypos - lastY) * forward * 0.01f; - camchanged = true; - } - lastX = xpos; - lastY = ypos; + if (xpos == lastX || ypos == lastY) return; // otherwise, clicking back into window causes re-start + if (leftMousePressed) { + // compute new camera parameters + phi -= (xpos - lastX) / width; + theta -= (ypos - lastY) / height; + theta = std::fmax(0.001f, std::fmin(theta, PI)); + camchanged = true; + } + else if (rightMousePressed) { + zoom += (ypos - lastY) / height; + zoom = std::fmax(0.1f, zoom); + camchanged = true; + } + else if (middleMousePressed) { + renderState = &scene->state; + Camera &cam = renderState->camera; + glm::vec3 forward = cam.view; + forward.y = 0.0f; + forward = glm::normalize(forward); + glm::vec3 right = cam.right; + right.y = 0.0f; + right = glm::normalize(right); + + cam.lookAt -= (float)(xpos - lastX) * right * 0.01f; + cam.lookAt += (float)(ypos - lastY) * forward * 0.01f; + camchanged = true; + } + lastX = xpos; + lastY = ypos; } diff --git a/src/main.h b/src/main.h index fdb7d5d..5caf9f2 100644 --- a/src/main.h +++ b/src/main.h @@ -20,7 +20,7 @@ #include "utilities.h" #include "scene.h" -using namespace std; +//using namespace std; //------------------------------- //----------PATH TRACER---------- diff --git a/src/pathtrace.cu b/src/pathtrace.cu index c1ec122..52d1c36 100644 --- a/src/pathtrace.cu +++ b/src/pathtrace.cu @@ -4,6 +4,9 @@ #include #include #include +#include +#include +#include #include "sceneStructs.h" #include "scene.h" @@ -14,57 +17,63 @@ #include "intersections.h" #include "interactions.h" + +#define SORT_BY_MATERIAL 0 +#define CACHE_FIRST_INTERSECTION 0 +#define STREAM_COMPACTION 0 +#define DOF false +#define MOTION_BLUR 1 #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(); - cudaError_t err = cudaGetLastError(); - if (cudaSuccess == err) { - return; - } - - fprintf(stderr, "CUDA error"); - if (file) { - fprintf(stderr, " (%s:%d)", file, line); - } - fprintf(stderr, ": %s: %s\n", msg, cudaGetErrorString(err)); + cudaDeviceSynchronize(); + cudaError_t err = cudaGetLastError(); + if (cudaSuccess == err) { + return; + } + + fprintf(stderr, "CUDA error"); + if (file) { + fprintf(stderr, " (%s:%d)", file, line); + } + fprintf(stderr, ": %s: %s\n", msg, cudaGetErrorString(err)); # ifdef _WIN32 - getchar(); + getchar(); # endif - exit(EXIT_FAILURE); + exit(EXIT_FAILURE); #endif } __host__ __device__ thrust::default_random_engine makeSeededRandomEngine(int iter, int index, int depth) { - int h = utilhash((1 << 31) | (depth << 22) | iter) ^ utilhash(index); - return thrust::default_random_engine(h); + int h = utilhash((1 << 31) | (depth << 22) | iter) ^ utilhash(index); + return thrust::default_random_engine(h); } //Kernel that writes the image to the OpenGL PBO directly. __global__ void sendImageToPBO(uchar4* pbo, glm::ivec2 resolution, - int iter, glm::vec3* image) { - int x = (blockIdx.x * blockDim.x) + threadIdx.x; - int y = (blockIdx.y * blockDim.y) + threadIdx.y; - - if (x < resolution.x && y < resolution.y) { - int index = x + (y * resolution.x); - glm::vec3 pix = image[index]; - - glm::ivec3 color; - color.x = glm::clamp((int) (pix.x / iter * 255.0), 0, 255); - color.y = glm::clamp((int) (pix.y / iter * 255.0), 0, 255); - color.z = glm::clamp((int) (pix.z / iter * 255.0), 0, 255); - - // Each thread writes one pixel location in the texture (textel) - pbo[index].w = 0; - pbo[index].x = color.x; - pbo[index].y = color.y; - pbo[index].z = color.z; - } + int iter, glm::vec3* image) { + int x = (blockIdx.x * blockDim.x) + threadIdx.x; + int y = (blockIdx.y * blockDim.y) + threadIdx.y; + + if (x < resolution.x && y < resolution.y) { + int index = x + (y * resolution.x); + glm::vec3 pix = image[index]; + + glm::ivec3 color; + color.x = glm::clamp((int)(pix.x / iter * 255.0), 0, 255); + color.y = glm::clamp((int)(pix.y / iter * 255.0), 0, 255); + color.z = glm::clamp((int)(pix.z / iter * 255.0), 0, 255); + + // Each thread writes one pixel location in the texture (textel) + pbo[index].w = 0; + pbo[index].x = color.x; + pbo[index].y = color.y; + pbo[index].z = color.z; + } } static Scene * hst_scene = NULL; @@ -73,42 +82,94 @@ static Geom * dev_geoms = NULL; static Material * dev_materials = NULL; static PathSegment * dev_paths = NULL; static ShadeableIntersection * dev_intersections = NULL; +static ShadeableIntersection * dev_first_intersections = NULL; + +bool cache_first_intersection = false; + // TODO: static variables for device memory, any extra info you need, etc // ... void pathtraceInit(Scene *scene) { - hst_scene = scene; - const Camera &cam = hst_scene->state.camera; - const int pixelcount = cam.resolution.x * cam.resolution.y; + hst_scene = scene; + const Camera &cam = hst_scene->state.camera; + const int pixelcount = cam.resolution.x * cam.resolution.y; - cudaMalloc(&dev_image, pixelcount * sizeof(glm::vec3)); - cudaMemset(dev_image, 0, pixelcount * sizeof(glm::vec3)); + cudaMalloc(&dev_image, pixelcount * sizeof(glm::vec3)); + cudaMemset(dev_image, 0, pixelcount * sizeof(glm::vec3)); - cudaMalloc(&dev_paths, pixelcount * sizeof(PathSegment)); + 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); + cudaMalloc(&dev_geoms, scene->geoms.size() * sizeof(Geom)); + cudaMemcpy(dev_geoms, scene->geoms.data(), scene->geoms.size() * sizeof(Geom), cudaMemcpyHostToDevice); - cudaMalloc(&dev_materials, scene->materials.size() * sizeof(Material)); - cudaMemcpy(dev_materials, scene->materials.data(), scene->materials.size() * sizeof(Material), cudaMemcpyHostToDevice); + cudaMalloc(&dev_materials, scene->materials.size() * sizeof(Material)); + cudaMemcpy(dev_materials, scene->materials.data(), scene->materials.size() * sizeof(Material), cudaMemcpyHostToDevice); + + cudaMalloc(&dev_intersections, pixelcount * sizeof(ShadeableIntersection)); + cudaMemset(dev_intersections, 0, pixelcount * sizeof(ShadeableIntersection)); - cudaMalloc(&dev_intersections, pixelcount * sizeof(ShadeableIntersection)); - cudaMemset(dev_intersections, 0, pixelcount * sizeof(ShadeableIntersection)); + // TODO: initialize any extra device memeory you need + cache_first_intersection = false; - // TODO: initialize any extra device memeory you need + cudaMalloc(&dev_first_intersections, pixelcount * sizeof(ShadeableIntersection)); + cudaMemset(dev_first_intersections, 0, pixelcount * sizeof(ShadeableIntersection)); - checkCUDAError("pathtraceInit"); + checkCUDAError("pathtraceInit"); } void pathtraceFree() { - cudaFree(dev_image); // no-op if dev_image is null - cudaFree(dev_paths); - cudaFree(dev_geoms); - cudaFree(dev_materials); - cudaFree(dev_intersections); - // TODO: clean up any extra device memory you created - - checkCUDAError("pathtraceFree"); + cudaFree(dev_image); // no-op if dev_image is null + cudaFree(dev_paths); + cudaFree(dev_geoms); + cudaFree(dev_materials); + cudaFree(dev_intersections); + // TODO: clean up any extra device memory you created + cudaFree(dev_first_intersections); + checkCUDAError("pathtraceFree"); +} + + +//Reference: PBRT source code +//ConcentricSampleDisk +__device__ glm::vec2 ConcentricSampleDisk(float u1, float u2) { + float r, theta; + float a, b; + // Map uniform random numbers to $[-1,1]^2$ + float sx = 2 * u1 - 1; + float sy = 2 * u2 - 1; + if (sx == 0.0 && sy == 0.0) { + return glm::vec2(0.f); + } + if (sx >= -sy) { + if (sx > sy) { + // Handle first region of disk + r = sx; + if (sy > 0.0) theta = sy / r; + else theta = 8.0f + sy / r; + } + else { + // Handle second region of disk + r = sy; + theta = 2.0f - sx / r; + } + } + else { + if (sx <= sy) { + // Handle third region of disk + r = -sx; + theta = 4.0f - sy / r; + } + else { + // Handle fourth region of disk + r = -sy; + theta = 6.0f + sx / r; + } + } + theta *= PI / 4.f; + a = r * cosf(theta); + b = r * sinf(theta); + glm::vec2 returnValue(a, b); + return returnValue; } /** @@ -119,24 +180,49 @@ void pathtraceFree() { * motion blur - jitter rays "in time" * lens effect - jitter ray origin positions based on a lens */ +//__global__ void generateRayFromCamera(Camera cam, int iter, int traceDepth, PathSegment* pathSegments) __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; + thrust::default_random_engine rng = makeSeededRandomEngine(iter, (x + (y * cam.resolution.x)), 0); + thrust::uniform_real_distribution u01(0, 1); + if (x < cam.resolution.x && y < cam.resolution.y) { + //Calculate index 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); + + // Depth of field + // pbrt page 313 - 318 + // The range of distances from the lens at which objects appear in focus is called the len's depth of field // 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 + u01(rng)) - (float)cam.resolution.x * 0.5f) + - cam.up * cam.pixelLength.y * ((float)(y + u01(rng)) - (float)cam.resolution.y * 0.5f) ); + //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) + // ); + if (DOF == true) { + glm::vec2 returnValue = ConcentricSampleDisk(u01(rng), u01(rng)); + returnValue.x *= cam.lensRadius; + returnValue.y *= cam.lensRadius; + //pbrt page 318 + float ft = glm::abs(cam.focalDistance / cam.view.z); + glm::vec3 Pfocus = segment.ray.direction * ft + segment.ray.origin; + + segment.ray.origin += returnValue.x * cam.right + returnValue.y * cam.up; + segment.ray.direction = glm::normalize(Pfocus - segment.ray.origin); + } + segment.pixelIndex = index; segment.remainingBounces = traceDepth; } @@ -156,11 +242,13 @@ __global__ void computeIntersections( ) { int path_index = blockIdx.x * blockDim.x + threadIdx.x; - if (path_index < num_paths) { - PathSegment pathSegment = pathSegments[path_index]; + PathSegment pathSegment = pathSegments[path_index]; + if (pathSegment.remainingBounces < 0) { + return; + } float t; glm::vec3 intersect_point; glm::vec3 normal; @@ -201,6 +289,7 @@ __global__ void computeIntersections( if (hit_geom_index == -1) { intersections[path_index].t = -1.0f; + intersections[path_index].materialId = -1; } else { @@ -208,6 +297,7 @@ __global__ void computeIntersections( intersections[path_index].t = t_min; intersections[path_index].materialId = geoms[hit_geom_index].materialid; intersections[path_index].surfaceNormal = normal; + intersections[path_index].position = intersect_point; } } } @@ -221,48 +311,59 @@ __global__ void computeIntersections( // Note that this shader does NOT do a BSDF evaluation! // Your shaders should handle that - this can allow techniques such as // bump mapping. -__global__ void shadeFakeMaterial ( - int iter - , int num_paths +__global__ void shadeFakeMaterial( + 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) { // 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); - } - // 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 - } - // 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); - } - } + int idx = blockIdx.x * blockDim.x + threadIdx.x; + if (idx < num_paths) + { + if (pathSegments[idx].remainingBounces < 0) { + return; + } + // Set up the RNG + // LOOK: this is how you use thrust's RNG! Please look at + // makeSeededRandomEngine as well. + ShadeableIntersection intersection = shadeableIntersections[idx]; + if (intersection.t > 0.0f) { // if the intersection exists... + thrust::default_random_engine rng = makeSeededRandomEngine(iter, idx, pathSegments[idx].remainingBounces); + 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 = -1; + } + // 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], intersection.position, intersection.surfaceNormal, material, rng); + if (pathSegments[idx].remainingBounces < 0) { + pathSegments[idx].color = glm::vec3(0.f); + } + //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].remainingBounces = -1; + } + } } // Add the current iteration's output to the overall image @@ -270,124 +371,200 @@ __global__ void finalGather(int nPaths, glm::vec3 * image, PathSegment * iterati { int index = (blockIdx.x * blockDim.x) + threadIdx.x; - if (index < nPaths) + if (index < nPaths && iterationPaths[index].remainingBounces < 0) { PathSegment iterationPath = iterationPaths[index]; image[iterationPath.pixelIndex] += iterationPath.color; } } +//thrust document +struct isPathAlive +{ + __host__ __device__ bool operator()(const PathSegment& path_segment) + { + return path_segment.remainingBounces >= 0; + } +}; + +//Sort by material +struct MaterialComparator +{ + __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 - */ +* 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) { - const int traceDepth = hst_scene->state.traceDepth; - const Camera &cam = hst_scene->state.camera; - const int pixelcount = cam.resolution.x * cam.resolution.y; + const int traceDepth = hst_scene->state.traceDepth; + const Camera &cam = hst_scene->state.camera; + const int pixelcount = cam.resolution.x * cam.resolution.y; + const Geom *hstSceneGeoms = &(hst_scene->geoms)[0]; + Geom *motionBlurGeoms = &(hst_scene->geoms)[0]; // 2D block for generating ray from camera - const dim3 blockSize2d(8, 8); - const dim3 blocksPerGrid2d( - (cam.resolution.x + blockSize2d.x - 1) / blockSize2d.x, - (cam.resolution.y + blockSize2d.y - 1) / blockSize2d.y); + const dim3 blockSize2d(8, 8); + const dim3 blocksPerGrid2d( + (cam.resolution.x + blockSize2d.x - 1) / blockSize2d.x, + (cam.resolution.y + blockSize2d.y - 1) / blockSize2d.y); // 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); + /////////////////////////////////////////////////////////////////////////// + + // 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 + + //toggle between motion blur and non-motion-blur +#if MOTION_BLUR + thrust::default_random_engine rng = makeSeededRandomEngine(iter, hst_scene->geoms.size(), traceDepth); + thrust::uniform_real_distribution u02PI(0, TWO_PI); + + for (int i = 0; i < hst_scene->geoms.size(); i++) { + motionBlurGeoms[i] = hstSceneGeoms[i]; + motionBlurGeoms[i].translation.x += hstSceneGeoms[i].motion.x * 0.08 * cos(u02PI(rng)); + motionBlurGeoms[i].translation.y += hstSceneGeoms[i].motion.y * 0.08 * cos(u02PI(rng)); + motionBlurGeoms[i].translation.z += hstSceneGeoms[i].motion.z * 0.08 * cos(u02PI(rng)); + //keep on updating all its transform information + motionBlurGeoms[i].transform = utilityCore::buildTransformationMatrix(motionBlurGeoms[i].translation, motionBlurGeoms[i].rotation, motionBlurGeoms[i].scale); + motionBlurGeoms[i].invTranspose = glm::inverseTranspose(motionBlurGeoms[i].transform); + motionBlurGeoms[i].inverseTransform = glm::inverse(motionBlurGeoms[i].transform); + } + + cudaMemcpy(dev_geoms, motionBlurGeoms, hst_scene->geoms.size() * sizeof(Geom), cudaMemcpyHostToDevice); + +#else + cudaMemcpy(dev_geoms, motionBlurGeoms, hst_scene->geoms.size() * sizeof(Geom), cudaMemcpyHostToDevice); + +#endif + + // generateRayFromCamera <<>>(cam, iter, traceDepth, dev_paths); + generateRayFromCamera << > >(cam, iter, traceDepth, dev_paths); checkCUDAError("generate camera ray"); int depth = 0; + thrust::device_ptr thrust_dev_paths(dev_paths); PathSegment* dev_path_end = dev_paths + pixelcount; int num_paths = dev_path_end - dev_paths; - + int num_paths_alive = num_paths; // --- PathSegment Tracing Stage --- // Shoot ray into scene, bounce between objects, push shading chunks - bool iterationComplete = false; + bool iterationComplete = false; + + //*******************************************************// + //******************big while begin**********************// while (!iterationComplete) { + // clean shading chunks + cudaMemset(dev_intersections, 0, pixelcount * sizeof(ShadeableIntersection)); + // tracing + dim3 numblocksPathSegmentTracing = (num_paths_alive + blockSize1d - 1) / blockSize1d; + if (depth == 0 && cache_first_intersection) { + cudaMemcpy(dev_intersections, dev_first_intersections, num_paths * sizeof(dev_first_intersections[0]), cudaMemcpyDeviceToDevice); + } + else { + computeIntersections << > > ( + depth + , num_paths_alive + , dev_paths + , dev_geoms + , hst_scene->geoms.size() + , dev_intersections + ); + } - // clean shading chunks - cudaMemset(dev_intersections, 0, pixelcount * sizeof(ShadeableIntersection)); + checkCUDAError("trace one bounce"); + cudaDeviceSynchronize(); + depth++; - // 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. - } + //Toggle between cache first intersection +#if CACHE_FIRST_INTERSECTION + if (!cache_first_intersection) { + cache_first_intersection = true; + cudaMemcpy(dev_first_intersections, dev_intersections, num_paths * sizeof(dev_first_intersections[0]), cudaMemcpyDeviceToDevice); + } +#endif - // Assemble this iteration and apply it to the image - dim3 numBlocksPixels = (pixelcount + blockSize1d - 1) / blockSize1d; - finalGather<<>>(num_paths, dev_image, dev_paths); + //Toggle sort by material. WHY SO SLOW......(not using it) +#if SORT_BY_MATERIAL + //Sort by material + thrust::sort_by_key(thrust::device, dev_intersections, dev_intersections + num_paths_alive, dev_paths, MaterialComparator()); +#endif + + // 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_alive, + dev_intersections, + dev_paths, + dev_materials + ); + + //Toggle stream compaction + //Google group reminder: using thrust::partition +#if STREAM_COMPACTION + num_paths_alive = thrust::partition(thrust::device, thrust_dev_paths, thrust_dev_paths + num_paths_alive, isPathAlive()) - thrust_dev_paths; + iterationComplete = (num_paths_alive <= 0); // TODO: should be based off stream compaction results. +#else + num_paths_alive = num_paths; + iterationComplete = (num_paths_alive <= 0) || depth > traceDepth; +#endif + + } + //*******************************************************// + //*****************big while end*************************// + // Assemble this iteration and apply it to the image + dim3 numBlocksPixels = (pixelcount + blockSize1d - 1) / blockSize1d; + finalGather << > >(num_paths, dev_image, dev_paths); - /////////////////////////////////////////////////////////////////////////// + /////////////////////////////////////////////////////////////////////////// - // Send results to OpenGL buffer for rendering - sendImageToPBO<<>>(pbo, cam.resolution, iter, dev_image); + // Send results to OpenGL buffer for rendering + sendImageToPBO << > >(pbo, cam.resolution, iter, dev_image); - // Retrieve image from GPU - cudaMemcpy(hst_scene->state.image.data(), dev_image, - pixelcount * sizeof(glm::vec3), cudaMemcpyDeviceToHost); + // Retrieve image from GPU + cudaMemcpy(hst_scene->state.image.data(), dev_image, + pixelcount * sizeof(glm::vec3), cudaMemcpyDeviceToHost); - checkCUDAError("pathtrace"); + checkCUDAError("pathtrace"); } diff --git a/src/scene.cpp b/src/scene.cpp index cbae043..e866bd5 100644 --- a/src/scene.cpp +++ b/src/scene.cpp @@ -74,8 +74,9 @@ 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(), "MOTION") == 0) { + newGeom.motion = glm::vec3(atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str())); + } utilityCore::safeGetline(fp_in, line); } @@ -114,6 +115,9 @@ int Scene::loadCamera() { } } + // initialize + camera.focalDistance = 0.f; + camera.lensRadius = 0.f; string line; utilityCore::safeGetline(fp_in, line); while (!line.empty() && fp_in.good()) { @@ -124,7 +128,11 @@ 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())); - } + } else if (strcmp(tokens[0].c_str(), "LENSRADIUS") == 0) { + camera.lensRadius = atof(tokens[1].c_str()); + } else if (strcmp(tokens[0].c_str(), "FOCALDISTANCE") == 0) { + camera.focalDistance = atof(tokens[1].c_str()); + } utilityCore::safeGetline(fp_in, line); } diff --git a/src/sceneStructs.h b/src/sceneStructs.h index b38b820..235a1b5 100644 --- a/src/sceneStructs.h +++ b/src/sceneStructs.h @@ -26,6 +26,7 @@ struct Geom { glm::mat4 transform; glm::mat4 inverseTransform; glm::mat4 invTranspose; + glm::vec3 motion; //for motion blur }; struct Material { @@ -49,6 +50,8 @@ struct Camera { glm::vec3 right; glm::vec2 fov; glm::vec2 pixelLength; + float lensRadius; + float focalDistance; }; struct RenderState { @@ -64,7 +67,7 @@ struct PathSegment { glm::vec3 color; int pixelIndex; int remainingBounces; -}; + }; // Use with a corresponding PathSegment to do: // 1) color contribution computation @@ -73,4 +76,5 @@ struct ShadeableIntersection { float t; glm::vec3 surfaceNormal; int materialId; + glm::vec3 position; //added information }; diff --git a/stream_compaction/CMakeLists.txt b/stream_compaction/CMakeLists.txt index ac358c9..d525972 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_30 )