diff --git a/README.md b/README.md index 110697c..c204fff 100644 --- a/README.md +++ b/README.md @@ -3,11 +3,55 @@ 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) +* Alexander Perry +* Tested on: Windows 10, i5-2410M @ 2.30GHz 8GB, NVS 4200M (personal computer) -### (TODO: Your README) +![](./build/.2017-10-01_22-05-47z.325samp.png) -*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. +This is my most representative image of the ray tracer. It includes antialiasing and path compression. It took a couple minutes to render. +Due to the low stats of my GPU I was unable to get high occupancy on many of the kernels. Several take a couple dozen registers to run with not much room to optimization. This limited the occupancy in the long running kernels to 33-46% and limited the overall performance of the renderer. +## Depth of field + +![](./build/.2017-10-02_19-39-46z.622samp.png) + +![](./build/.2017-10-02_19-45-37z.709samp.png) + +These images showcase the depth of field effect. The first image has the focal distance set to near the center sphere while the second has the focal distance set to the distance of the rear wall. In the second image, it is clear to see the out-of-focusness increase along the wall toward the camera as the bleed-over increases between the walls and the floor/ceiling. + +Depth of field was implemented by jittering ray origins in a disc. + +![](./build/.2017-10-02_19-35-20z.20samp.png) + +![](./build/.2017-10-02_19-35-49z.40samp.png) + +![](./build/.2017-10-02_19-36-23z.120samp.png) + +This is a series of images with reducing lensRadius. The focal distance is closer than everything in the scene but one can see how the radius of the lens affects the blur. + +## Antialias + +![](./img/no_aa.png) +![](./img/with_aa.png) + +These images were created with no antialias for the first image and with antialiasing on the second. Both images were approximately 325 samples. Through the noise, it is possible to tell that the image on the left is a bit blockier along the edge. + +Antialiasing was accomplished by jittering rays in a pixel. This is possible to run at the same time as depth of field. Antialiasing runs first, followed by depth of field in that case. + +## Direct lighting + +Direct lighting was accomplished by sending a feeler ray to a random light source of all rays still alive at the end of the iterations. This helps the image converge faster as more pixels have a light source in them at the end. + +## Bloopers + +![](./build/.2017-09-29_00-13-42z.22samp.png) + +This was caused by drawing all pixels at the end rather than only valid pixels. Due to stream compaction, pixels at the end of the array (at the bottom of the screen) showed up multiple times in the array and were significantly brighter. + +![](./build/.2017-09-29_01-50-03z.494samp.png) + +This is standard "shadow acne". Due to floating point error, intersections were occuring in the wall rather than in front of it. + +![](./build/.2017-10-01_03-30-23z.1samp.png) + +This was early in setting up direct lighting. diff --git a/build/.2017-09-28_20-54-46z.1036samp.png b/build/.2017-09-28_20-54-46z.1036samp.png new file mode 100644 index 0000000..acebe35 Binary files /dev/null and b/build/.2017-09-28_20-54-46z.1036samp.png differ diff --git a/build/.2017-09-28_21-07-32z.13samp.png b/build/.2017-09-28_21-07-32z.13samp.png new file mode 100644 index 0000000..add2a52 Binary files /dev/null and b/build/.2017-09-28_21-07-32z.13samp.png differ diff --git a/build/.2017-09-28_21-29-26z.129samp.png b/build/.2017-09-28_21-29-26z.129samp.png new file mode 100644 index 0000000..65ce19f Binary files /dev/null and b/build/.2017-09-28_21-29-26z.129samp.png differ diff --git a/build/.2017-09-28_21-38-06z.174samp.png b/build/.2017-09-28_21-38-06z.174samp.png new file mode 100644 index 0000000..3e8e24f Binary files /dev/null and b/build/.2017-09-28_21-38-06z.174samp.png differ diff --git a/build/.2017-09-28_21-43-39z.32samp.png b/build/.2017-09-28_21-43-39z.32samp.png new file mode 100644 index 0000000..2bab966 Binary files /dev/null and b/build/.2017-09-28_21-43-39z.32samp.png differ diff --git a/build/.2017-09-28_21-57-59z.6samp.png b/build/.2017-09-28_21-57-59z.6samp.png new file mode 100644 index 0000000..b258189 Binary files /dev/null and b/build/.2017-09-28_21-57-59z.6samp.png differ diff --git a/build/.2017-09-28_22-00-02z.3samp.png b/build/.2017-09-28_22-00-02z.3samp.png new file mode 100644 index 0000000..376a279 Binary files /dev/null and b/build/.2017-09-28_22-00-02z.3samp.png differ diff --git a/build/.2017-09-28_22-01-21z.3samp.png b/build/.2017-09-28_22-01-21z.3samp.png new file mode 100644 index 0000000..11499f6 Binary files /dev/null and b/build/.2017-09-28_22-01-21z.3samp.png differ diff --git a/build/.2017-09-28_22-06-03z.37samp.png b/build/.2017-09-28_22-06-03z.37samp.png new file mode 100644 index 0000000..c929908 Binary files /dev/null and b/build/.2017-09-28_22-06-03z.37samp.png differ diff --git a/build/.2017-09-28_22-11-11z.33samp.png b/build/.2017-09-28_22-11-11z.33samp.png new file mode 100644 index 0000000..500c685 Binary files /dev/null and b/build/.2017-09-28_22-11-11z.33samp.png differ diff --git a/build/.2017-09-28_22-14-23z.35samp.png b/build/.2017-09-28_22-14-23z.35samp.png new file mode 100644 index 0000000..dbf99bf Binary files /dev/null and b/build/.2017-09-28_22-14-23z.35samp.png differ diff --git a/build/.2017-09-28_22-18-07z.12samp.png b/build/.2017-09-28_22-18-07z.12samp.png new file mode 100644 index 0000000..b718fcd Binary files /dev/null and b/build/.2017-09-28_22-18-07z.12samp.png differ diff --git a/build/.2017-09-28_22-19-07z.7samp.png b/build/.2017-09-28_22-19-07z.7samp.png new file mode 100644 index 0000000..6cc6af2 Binary files /dev/null and b/build/.2017-09-28_22-19-07z.7samp.png differ diff --git a/build/.2017-09-28_22-19-53z.80samp.png b/build/.2017-09-28_22-19-53z.80samp.png new file mode 100644 index 0000000..45faf3c Binary files /dev/null and b/build/.2017-09-28_22-19-53z.80samp.png differ diff --git a/build/.2017-09-28_22-32-12z.3samp.png b/build/.2017-09-28_22-32-12z.3samp.png new file mode 100644 index 0000000..69aa95e Binary files /dev/null and b/build/.2017-09-28_22-32-12z.3samp.png differ diff --git a/build/.2017-09-28_22-45-19z.95samp.png b/build/.2017-09-28_22-45-19z.95samp.png new file mode 100644 index 0000000..d41a222 Binary files /dev/null and b/build/.2017-09-28_22-45-19z.95samp.png differ diff --git a/build/.2017-09-29_00-03-20z.12samp.png b/build/.2017-09-29_00-03-20z.12samp.png new file mode 100644 index 0000000..019ae8f Binary files /dev/null and b/build/.2017-09-29_00-03-20z.12samp.png differ diff --git a/build/.2017-09-29_00-13-42z.22samp.png b/build/.2017-09-29_00-13-42z.22samp.png new file mode 100644 index 0000000..5d3abf3 Binary files /dev/null and b/build/.2017-09-29_00-13-42z.22samp.png differ diff --git a/build/.2017-09-29_00-21-30z.636samp.png b/build/.2017-09-29_00-21-30z.636samp.png new file mode 100644 index 0000000..81adb1a Binary files /dev/null and b/build/.2017-09-29_00-21-30z.636samp.png differ diff --git a/build/.2017-09-29_00-28-33z.397samp.png b/build/.2017-09-29_00-28-33z.397samp.png new file mode 100644 index 0000000..b552ecb Binary files /dev/null and b/build/.2017-09-29_00-28-33z.397samp.png differ diff --git a/build/.2017-09-29_00-28-33z.566samp.png b/build/.2017-09-29_00-28-33z.566samp.png new file mode 100644 index 0000000..fd037f4 Binary files /dev/null and b/build/.2017-09-29_00-28-33z.566samp.png differ diff --git a/build/.2017-09-29_01-47-02z.34samp.png b/build/.2017-09-29_01-47-02z.34samp.png new file mode 100644 index 0000000..c88a2fb Binary files /dev/null and b/build/.2017-09-29_01-47-02z.34samp.png differ diff --git a/build/.2017-09-29_01-47-02z.78samp.png b/build/.2017-09-29_01-47-02z.78samp.png new file mode 100644 index 0000000..a439a89 Binary files /dev/null and b/build/.2017-09-29_01-47-02z.78samp.png differ diff --git a/build/.2017-09-29_01-50-03z.160samp.png b/build/.2017-09-29_01-50-03z.160samp.png new file mode 100644 index 0000000..93d37c0 Binary files /dev/null and b/build/.2017-09-29_01-50-03z.160samp.png differ diff --git a/build/.2017-09-29_01-50-03z.491samp.png b/build/.2017-09-29_01-50-03z.491samp.png new file mode 100644 index 0000000..b3103e6 Binary files /dev/null and b/build/.2017-09-29_01-50-03z.491samp.png differ diff --git a/build/.2017-09-29_01-50-03z.494samp.png b/build/.2017-09-29_01-50-03z.494samp.png new file mode 100644 index 0000000..c60c9bb Binary files /dev/null and b/build/.2017-09-29_01-50-03z.494samp.png differ diff --git a/build/.2017-09-29_01-57-48z.675samp.png b/build/.2017-09-29_01-57-48z.675samp.png new file mode 100644 index 0000000..9c8a4b9 Binary files /dev/null and b/build/.2017-09-29_01-57-48z.675samp.png differ diff --git a/build/.2017-09-30_18-52-05z.341samp.png b/build/.2017-09-30_18-52-05z.341samp.png new file mode 100644 index 0000000..9ac682f Binary files /dev/null and b/build/.2017-09-30_18-52-05z.341samp.png differ diff --git a/build/.2017-09-30_19-03-03z.40samp.png b/build/.2017-09-30_19-03-03z.40samp.png new file mode 100644 index 0000000..51550e9 Binary files /dev/null and b/build/.2017-09-30_19-03-03z.40samp.png differ diff --git a/build/.2017-09-30_19-04-35z.11samp.png b/build/.2017-09-30_19-04-35z.11samp.png new file mode 100644 index 0000000..2ea5b56 Binary files /dev/null and b/build/.2017-09-30_19-04-35z.11samp.png differ diff --git a/build/.2017-09-30_20-56-25z.11samp.png b/build/.2017-09-30_20-56-25z.11samp.png new file mode 100644 index 0000000..2ea5b56 Binary files /dev/null and b/build/.2017-09-30_20-56-25z.11samp.png differ diff --git a/build/.2017-09-30_21-23-11z.11samp.png b/build/.2017-09-30_21-23-11z.11samp.png new file mode 100644 index 0000000..2ea5b56 Binary files /dev/null and b/build/.2017-09-30_21-23-11z.11samp.png differ diff --git a/build/.2017-09-30_21-29-44z.7samp.png b/build/.2017-09-30_21-29-44z.7samp.png new file mode 100644 index 0000000..75a6caa Binary files /dev/null and b/build/.2017-09-30_21-29-44z.7samp.png differ diff --git a/build/.2017-09-30_21-30-26z.3samp.png b/build/.2017-09-30_21-30-26z.3samp.png new file mode 100644 index 0000000..63f3384 Binary files /dev/null and b/build/.2017-09-30_21-30-26z.3samp.png differ diff --git a/build/.2017-09-30_21-36-36z.4samp.png b/build/.2017-09-30_21-36-36z.4samp.png new file mode 100644 index 0000000..9d15871 Binary files /dev/null and b/build/.2017-09-30_21-36-36z.4samp.png differ diff --git a/build/.2017-09-30_21-39-00z.6samp.png b/build/.2017-09-30_21-39-00z.6samp.png new file mode 100644 index 0000000..35ee1ec Binary files /dev/null and b/build/.2017-09-30_21-39-00z.6samp.png differ diff --git a/build/.2017-09-30_21-55-40z.41samp.png b/build/.2017-09-30_21-55-40z.41samp.png new file mode 100644 index 0000000..31da407 Binary files /dev/null and b/build/.2017-09-30_21-55-40z.41samp.png differ diff --git a/build/.2017-09-30_21-58-09z.15samp.png b/build/.2017-09-30_21-58-09z.15samp.png new file mode 100644 index 0000000..5b92b46 Binary files /dev/null and b/build/.2017-09-30_21-58-09z.15samp.png differ diff --git a/build/.2017-09-30_23-58-31z.12samp.png b/build/.2017-09-30_23-58-31z.12samp.png new file mode 100644 index 0000000..c64a413 Binary files /dev/null and b/build/.2017-09-30_23-58-31z.12samp.png differ diff --git a/build/.2017-10-01_00-01-19z.8samp.png b/build/.2017-10-01_00-01-19z.8samp.png new file mode 100644 index 0000000..cc47a4d Binary files /dev/null and b/build/.2017-10-01_00-01-19z.8samp.png differ diff --git a/build/.2017-10-01_00-03-11z.5samp.png b/build/.2017-10-01_00-03-11z.5samp.png new file mode 100644 index 0000000..4abded9 Binary files /dev/null and b/build/.2017-10-01_00-03-11z.5samp.png differ diff --git a/build/.2017-10-01_00-07-24z.4samp.png b/build/.2017-10-01_00-07-24z.4samp.png new file mode 100644 index 0000000..962e3a5 Binary files /dev/null and b/build/.2017-10-01_00-07-24z.4samp.png differ diff --git a/build/.2017-10-01_01-04-33z.2samp.png b/build/.2017-10-01_01-04-33z.2samp.png new file mode 100644 index 0000000..5c4c434 Binary files /dev/null and b/build/.2017-10-01_01-04-33z.2samp.png differ diff --git a/build/.2017-10-01_01-06-13z.1samp.png b/build/.2017-10-01_01-06-13z.1samp.png new file mode 100644 index 0000000..173d165 Binary files /dev/null and b/build/.2017-10-01_01-06-13z.1samp.png differ diff --git a/build/.2017-10-01_01-39-47z.9samp.png b/build/.2017-10-01_01-39-47z.9samp.png new file mode 100644 index 0000000..67a642e Binary files /dev/null and b/build/.2017-10-01_01-39-47z.9samp.png differ diff --git a/build/.2017-10-01_01-40-53z.5samp.png b/build/.2017-10-01_01-40-53z.5samp.png new file mode 100644 index 0000000..4abded9 Binary files /dev/null and b/build/.2017-10-01_01-40-53z.5samp.png differ diff --git a/build/.2017-10-01_01-41-35z.2samp.png b/build/.2017-10-01_01-41-35z.2samp.png new file mode 100644 index 0000000..5c4c434 Binary files /dev/null and b/build/.2017-10-01_01-41-35z.2samp.png differ diff --git a/build/.2017-10-01_02-44-12z.7samp.png b/build/.2017-10-01_02-44-12z.7samp.png new file mode 100644 index 0000000..c87d495 Binary files /dev/null and b/build/.2017-10-01_02-44-12z.7samp.png differ diff --git a/build/.2017-10-01_02-47-22z.7samp.png b/build/.2017-10-01_02-47-22z.7samp.png new file mode 100644 index 0000000..c87d495 Binary files /dev/null and b/build/.2017-10-01_02-47-22z.7samp.png differ diff --git a/build/.2017-10-01_02-48-08z.466samp.png b/build/.2017-10-01_02-48-08z.466samp.png new file mode 100644 index 0000000..52c099a Binary files /dev/null and b/build/.2017-10-01_02-48-08z.466samp.png differ diff --git a/build/.2017-10-01_03-30-23z.1samp.png b/build/.2017-10-01_03-30-23z.1samp.png new file mode 100644 index 0000000..0594907 Binary files /dev/null and b/build/.2017-10-01_03-30-23z.1samp.png differ diff --git a/build/.2017-10-01_03-32-11z.9samp.png b/build/.2017-10-01_03-32-11z.9samp.png new file mode 100644 index 0000000..67a642e Binary files /dev/null and b/build/.2017-10-01_03-32-11z.9samp.png differ diff --git a/build/.2017-10-01_04-29-01z.8samp.png b/build/.2017-10-01_04-29-01z.8samp.png new file mode 100644 index 0000000..cc47a4d Binary files /dev/null and b/build/.2017-10-01_04-29-01z.8samp.png differ diff --git a/build/.2017-10-01_04-30-08z.3samp.png b/build/.2017-10-01_04-30-08z.3samp.png new file mode 100644 index 0000000..fc832b0 Binary files /dev/null and b/build/.2017-10-01_04-30-08z.3samp.png differ diff --git a/build/.2017-10-01_04-41-49z.6samp.png b/build/.2017-10-01_04-41-49z.6samp.png new file mode 100644 index 0000000..2e14bbd Binary files /dev/null and b/build/.2017-10-01_04-41-49z.6samp.png differ diff --git a/build/.2017-10-01_18-33-45z.85samp.png b/build/.2017-10-01_18-33-45z.85samp.png new file mode 100644 index 0000000..cf0582d Binary files /dev/null and b/build/.2017-10-01_18-33-45z.85samp.png differ diff --git a/build/.2017-10-01_18-37-35z.39samp.png b/build/.2017-10-01_18-37-35z.39samp.png new file mode 100644 index 0000000..86ca223 Binary files /dev/null and b/build/.2017-10-01_18-37-35z.39samp.png differ diff --git a/build/.2017-10-01_19-14-06z.50samp.png b/build/.2017-10-01_19-14-06z.50samp.png new file mode 100644 index 0000000..38d146f Binary files /dev/null and b/build/.2017-10-01_19-14-06z.50samp.png differ diff --git a/build/.2017-10-01_20-03-51z.8samp.png b/build/.2017-10-01_20-03-51z.8samp.png new file mode 100644 index 0000000..105f0ea Binary files /dev/null and b/build/.2017-10-01_20-03-51z.8samp.png differ diff --git a/build/.2017-10-01_20-41-04z.146samp.png b/build/.2017-10-01_20-41-04z.146samp.png new file mode 100644 index 0000000..70b5e90 Binary files /dev/null and b/build/.2017-10-01_20-41-04z.146samp.png differ diff --git a/build/.2017-10-01_21-08-32z.120samp.png b/build/.2017-10-01_21-08-32z.120samp.png new file mode 100644 index 0000000..6cfdb70 Binary files /dev/null and b/build/.2017-10-01_21-08-32z.120samp.png differ diff --git a/build/.2017-10-01_21-10-21z.30samp.png b/build/.2017-10-01_21-10-21z.30samp.png new file mode 100644 index 0000000..145e567 Binary files /dev/null and b/build/.2017-10-01_21-10-21z.30samp.png differ diff --git a/build/.2017-10-01_21-11-26z.217samp.png b/build/.2017-10-01_21-11-26z.217samp.png new file mode 100644 index 0000000..fd927c7 Binary files /dev/null and b/build/.2017-10-01_21-11-26z.217samp.png differ diff --git a/build/.2017-10-01_21-25-23z.45samp.png b/build/.2017-10-01_21-25-23z.45samp.png new file mode 100644 index 0000000..4cf9ca9 Binary files /dev/null and b/build/.2017-10-01_21-25-23z.45samp.png differ diff --git a/build/.2017-10-01_21-25-23z.568samp.png b/build/.2017-10-01_21-25-23z.568samp.png new file mode 100644 index 0000000..975f7b4 Binary files /dev/null and b/build/.2017-10-01_21-25-23z.568samp.png differ diff --git a/build/.2017-10-01_21-32-18z.70samp.png b/build/.2017-10-01_21-32-18z.70samp.png new file mode 100644 index 0000000..04f988a Binary files /dev/null and b/build/.2017-10-01_21-32-18z.70samp.png differ diff --git a/build/.2017-10-01_21-59-34z.147samp.png b/build/.2017-10-01_21-59-34z.147samp.png new file mode 100644 index 0000000..9d34bc6 Binary files /dev/null and b/build/.2017-10-01_21-59-34z.147samp.png differ diff --git a/build/.2017-10-01_22-01-28z.404samp.png b/build/.2017-10-01_22-01-28z.404samp.png new file mode 100644 index 0000000..15a7a20 Binary files /dev/null and b/build/.2017-10-01_22-01-28z.404samp.png differ diff --git a/build/.2017-10-01_22-05-47z.325samp.png b/build/.2017-10-01_22-05-47z.325samp.png new file mode 100644 index 0000000..b5a0a8a Binary files /dev/null and b/build/.2017-10-01_22-05-47z.325samp.png differ diff --git a/build/.2017-10-01_22-25-30z.1samp.png b/build/.2017-10-01_22-25-30z.1samp.png new file mode 100644 index 0000000..771b6e0 Binary files /dev/null and b/build/.2017-10-01_22-25-30z.1samp.png differ diff --git a/build/.2017-10-02_19-35-20z.20samp.png b/build/.2017-10-02_19-35-20z.20samp.png new file mode 100644 index 0000000..2a1142f Binary files /dev/null and b/build/.2017-10-02_19-35-20z.20samp.png differ diff --git a/build/.2017-10-02_19-35-49z.40samp.png b/build/.2017-10-02_19-35-49z.40samp.png new file mode 100644 index 0000000..879756f Binary files /dev/null and b/build/.2017-10-02_19-35-49z.40samp.png differ diff --git a/build/.2017-10-02_19-36-23z.120samp.png b/build/.2017-10-02_19-36-23z.120samp.png new file mode 100644 index 0000000..5230b29 Binary files /dev/null and b/build/.2017-10-02_19-36-23z.120samp.png differ diff --git a/build/.2017-10-02_19-38-17z.138samp.png b/build/.2017-10-02_19-38-17z.138samp.png new file mode 100644 index 0000000..a1a2178 Binary files /dev/null and b/build/.2017-10-02_19-38-17z.138samp.png differ diff --git a/build/.2017-10-02_19-39-46z.622samp.png b/build/.2017-10-02_19-39-46z.622samp.png new file mode 100644 index 0000000..4599441 Binary files /dev/null and b/build/.2017-10-02_19-39-46z.622samp.png differ diff --git a/build/.2017-10-02_19-45-37z.709samp.png b/build/.2017-10-02_19-45-37z.709samp.png new file mode 100644 index 0000000..4ceaf68 Binary files /dev/null and b/build/.2017-10-02_19-45-37z.709samp.png differ diff --git a/build/.2017-10-02_19-45-37z.901samp.png b/build/.2017-10-02_19-45-37z.901samp.png new file mode 100644 index 0000000..7c04278 Binary files /dev/null and b/build/.2017-10-02_19-45-37z.901samp.png differ diff --git a/build/.2017-10-02_19-54-33z.102samp.png b/build/.2017-10-02_19-54-33z.102samp.png new file mode 100644 index 0000000..d8dbe2d Binary files /dev/null and b/build/.2017-10-02_19-54-33z.102samp.png differ diff --git a/build/.2017-10-02_19-54-33z.57samp.png b/build/.2017-10-02_19-54-33z.57samp.png new file mode 100644 index 0000000..20be225 Binary files /dev/null and b/build/.2017-10-02_19-54-33z.57samp.png differ diff --git a/build/.2017-10-03_01-18-22z.325samp.png b/build/.2017-10-03_01-18-22z.325samp.png new file mode 100644 index 0000000..6dbd342 Binary files /dev/null and b/build/.2017-10-03_01-18-22z.325samp.png differ diff --git a/build/.2017-10-03_01-21-45z.324samp.png b/build/.2017-10-03_01-21-45z.324samp.png new file mode 100644 index 0000000..2fa51aa Binary files /dev/null and b/build/.2017-10-03_01-21-45z.324samp.png differ diff --git a/build/Release/.2017-10-01_04-01-15z.2samp.png b/build/Release/.2017-10-01_04-01-15z.2samp.png new file mode 100644 index 0000000..5c4c434 Binary files /dev/null and b/build/Release/.2017-10-01_04-01-15z.2samp.png differ diff --git a/build/Release/.2017-10-01_04-01-28z.2samp.png b/build/Release/.2017-10-01_04-01-28z.2samp.png new file mode 100644 index 0000000..5c4c434 Binary files /dev/null and b/build/Release/.2017-10-01_04-01-28z.2samp.png differ diff --git a/build/Release/.2017-10-01_04-04-59z.2samp.png b/build/Release/.2017-10-01_04-04-59z.2samp.png new file mode 100644 index 0000000..5c4c434 Binary files /dev/null and b/build/Release/.2017-10-01_04-04-59z.2samp.png differ diff --git a/img/no_aa.png b/img/no_aa.png new file mode 100644 index 0000000..7eb80e7 Binary files /dev/null and b/img/no_aa.png differ diff --git a/img/with_aa.png b/img/with_aa.png new file mode 100644 index 0000000..a3dfd88 Binary files /dev/null and b/img/with_aa.png differ diff --git a/scenes/cornell.txt b/scenes/cornell.txt index 83ff820..d4cff46 100644 --- a/scenes/cornell.txt +++ b/scenes/cornell.txt @@ -51,6 +51,7 @@ EMITTANCE 0 // Camera CAMERA RES 800 800 +//RES 200 200 FOVY 45 ITERATIONS 5000 DEPTH 8 @@ -58,6 +59,8 @@ FILE cornell EYE 0.0 5 10.5 LOOKAT 0 5 0 UP 0 1 0 +FOCALLENGTH 15.5 +LENSRADIUS 1.5 // Ceiling light diff --git a/scenes/sphere.txt b/scenes/sphere.txt index a74b545..5077990 100644 --- a/scenes/sphere.txt +++ b/scenes/sphere.txt @@ -10,13 +10,14 @@ EMITTANCE 5 // Camera CAMERA -RES 800 800 +//RES 800 800 +RES 200 200 FOVY 45 ITERATIONS 5000 DEPTH 8 FILE sphere -EYE 0.0 5 10.5 -LOOKAT 0 5 0 +EYE 0.0 0 10.5 +LOOKAT 0 0 0 UP 0 1 0 // Sphere diff --git a/src/interactions.h b/src/interactions.h index 5ce3628..6f00449 100644 --- a/src/interactions.h +++ b/src/interactions.h @@ -67,13 +67,36 @@ glm::vec3 calculateRandomDirectionInHemisphere( * You may need to change the parameter list for your purposes! */ __host__ __device__ -void scatterRay( - PathSegment & pathSegment, - glm::vec3 intersect, - glm::vec3 normal, - const Material &m, - thrust::default_random_engine &rng) { +void scatterRayUniform( + 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. + glm::vec3 dir = calculateRandomDirectionInHemisphere(normal, rng); + pathSegment.ray.direction = dir; + pathSegment.ray.origin = intersect+dir*0.00001f; // avoid shadow acne + --pathSegment.remainingBounces; + //pathSegment.color = pathSegment.remainingBounces > 0 ? + // pathSegment.color * m.color : glm::vec3(0.0f); + pathSegment.color = pathSegment.color * m.color; +} + +__host__ __device__ +void scatterRaySpecular( + PathSegment &pathSegment, + glm::vec3 &intersect, + glm::vec3 &normal, + const Material &m, + thrust::default_random_engine &rng) { + glm::vec3 dir = glm::reflect(pathSegment.ray.direction, normal); + pathSegment.ray.direction = dir; + pathSegment.ray.origin = intersect+dir*0.00001f; // avoid shadow acne + --pathSegment.remainingBounces; + //pathSegment.color = pathSegment.remainingBounces > 0 ? + // pathSegment.color * m.specular.color : glm::vec3(0.0f); + pathSegment.color = pathSegment.color* m.specular.color; } diff --git a/src/intersections.h b/src/intersections.h index 6f23872..eefb1e4 100644 --- a/src/intersections.h +++ b/src/intersections.h @@ -45,8 +45,8 @@ __host__ __device__ glm::vec3 multiplyMV(glm::mat4 m, glm::vec4 v) { * @param outside Output param for whether the ray came from outside. * @return Ray parameter `t` value. -1 if no intersection. */ -__host__ __device__ float boxIntersectionTest(Geom box, Ray r, - glm::vec3 &intersectionPoint, glm::vec3 &normal, bool &outside) { +__host__ __device__ float boxIntersectionTest(Geom &box, Ray &r, + glm::vec3 &intersectionPoint, glm::vec3 &normal/*, bool &outside*/) { Ray q; q.origin = multiplyMV(box.inverseTransform, glm::vec4(r.origin , 1.0f)); q.direction = glm::normalize(multiplyMV(box.inverseTransform, glm::vec4(r.direction, 0.0f))); @@ -76,11 +76,11 @@ __host__ __device__ float boxIntersectionTest(Geom box, Ray r, } if (tmax >= tmin && tmax > 0) { - outside = true; + //outside = true; if (tmin <= 0) { tmin = tmax; tmin_n = tmax_n; - outside = false; + //outside = false; } intersectionPoint = multiplyMV(box.transform, glm::vec4(getPointOnRay(q, tmin), 1.0f)); normal = glm::normalize(multiplyMV(box.transform, glm::vec4(tmin_n, 0.0f))); @@ -99,8 +99,8 @@ __host__ __device__ float boxIntersectionTest(Geom box, Ray r, * @param outside Output param for whether the ray came from outside. * @return Ray parameter `t` value. -1 if no intersection. */ -__host__ __device__ float sphereIntersectionTest(Geom sphere, Ray r, - glm::vec3 &intersectionPoint, glm::vec3 &normal, bool &outside) { +__host__ __device__ float sphereIntersectionTest(Geom &sphere, Ray &r, + glm::vec3 &intersectionPoint, glm::vec3 &normal/*, bool &outside*/) { float radius = .5; glm::vec3 ro = multiplyMV(sphere.inverseTransform, glm::vec4(r.origin, 1.0f)); @@ -120,6 +120,7 @@ __host__ __device__ float sphereIntersectionTest(Geom sphere, Ray r, float firstTerm = -vDotDirection; float t1 = firstTerm + squareRoot; float t2 = firstTerm - squareRoot; + bool outside; float t = 0; if (t1 < 0 && t2 < 0) { diff --git a/src/pathtrace.cu b/src/pathtrace.cu index c1ec122..70ab280 100644 --- a/src/pathtrace.cu +++ b/src/pathtrace.cu @@ -1,6 +1,7 @@ #include #include #include +#include #include #include #include @@ -16,6 +17,11 @@ #define ERRORCHECK 1 +#define CACHE_FIRST_BOUNCE 0 +#define ANTIALIAS 1 // Not compatible with CACH_FIRST_BOUNCE +#define DOF 1 // compatible with ANTIALIAS +#define SORT_MATERIALS 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 +73,12 @@ __global__ void sendImageToPBO(uchar4* pbo, glm::ivec2 resolution, } } +// Typedefs for sorting +typedef thrust::device_ptr SegPtr; +typedef thrust::device_ptr IntersectionPtr; +typedef thrust::tuple PtrTuple; +typedef thrust::zip_iterator ZipIterator; + static Scene * hst_scene = NULL; static glm::vec3 * dev_image = NULL; static Geom * dev_geoms = NULL; @@ -75,29 +87,57 @@ static PathSegment * dev_paths = NULL; static ShadeableIntersection * dev_intersections = NULL; // TODO: static variables for device memory, any extra info you need, etc // ... +static SegPtr dev_thrust_paths; +static IntersectionPtr dev_thrust_intersections; +static int *dev_material_indices = nullptr; +static thrust::device_ptr dev_thrust_material_indices; +static ShadeableIntersection *dev_first_bounce = NULL; +static bool first_bounce_cached = false; +static int * dev_light_indices = NULL; +static int num_lights; 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 + cudaMalloc(&dev_first_bounce, pixelcount * sizeof(ShadeableIntersection)); + cudaMemset(dev_first_bounce, 0, pixelcount * sizeof(ShadeableIntersection)); - checkCUDAError("pathtraceInit"); + cudaMalloc(&dev_material_indices, pixelcount * sizeof(int)); + cudaMemset(dev_material_indices, -1, pixelcount * sizeof(int)); + + dev_thrust_paths = SegPtr(dev_paths); + dev_thrust_intersections = IntersectionPtr(dev_intersections); + dev_thrust_material_indices = thrust::device_ptr(dev_material_indices); + + // set up lights + cudaMalloc(&dev_light_indices, scene->geoms.size() * sizeof(Geom)); + int light_index = 0; + for (int i = 0; i < scene->geoms.size(); ++i) { + if (scene->materials[scene->geoms[i].materialid].emittance > 0.0f) { + //cudaMemcpy(&dev_light_indices[light_index], &dev_geoms[i], sizeof(Geom), cudaMemcpyDeviceToDevice); + cudaMemset(&dev_light_indices[light_index], i, sizeof(int)); + ++light_index; + } + num_lights = light_index; + } + + checkCUDAError("pathtraceInit"); } void pathtraceFree() { @@ -119,27 +159,73 @@ 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) -{ - int x = (blockIdx.x * blockDim.x) + threadIdx.x; - int y = (blockIdx.y * blockDim.y) + threadIdx.y; +__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; - if (x < cam.resolution.x && y < cam.resolution.y) { - int index = x + (y * cam.resolution.x); - PathSegment & segment = pathSegments[index]; + if (x < cam.resolution.x && y < cam.resolution.y) { + int index = x + (y * cam.resolution.x); + PathSegment & segment = pathSegments[index]; - segment.ray.origin = cam.position; + thrust::default_random_engine rng = makeSeededRandomEngine(iter, index, 0); + thrust::uniform_real_distribution u5(-0.5, 0.5); + thrust::uniform_real_distribution u01(0, 1); + + segment.ray.origin = cam.position; segment.color = glm::vec3(1.0f, 1.0f, 1.0f); - // 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) - ); + // TODO: implement antialiasing by jittering the ray +#if ANTIALIAS + segment.ray.direction = glm::normalize(cam.view + - cam.right * cam.pixelLength.x * ((float)x - (float)cam.resolution.x * 0.5f + u5(rng)) + - cam.up * cam.pixelLength.y * ((float)y - (float)cam.resolution.y * 0.5f + u5(rng))); +#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 DOF + float f_dist = cam.focalLength / glm::dot(segment.ray.direction, cam.view); + glm::vec3 focal_point = getPointOnRay(segment.ray, f_dist); + float t = u01(rng); + float s = u01(rng); + segment.ray.origin = segment.ray.origin + cam.right*sqrt(t)*cos(2*PI*s) * cam.lensDiameter + + cam.up *sqrt(t)*sin(2*PI*s) * cam.lensDiameter; + segment.ray.direction = glm::normalize(focal_point - segment.ray.origin); +#endif + + segment.pixelIndex = index; + segment.remainingBounces = traceDepth; + } +} - segment.pixelIndex = index; - segment.remainingBounces = traceDepth; - } +__device__ void FindNearestIntersection( + int geoms_size, Ray &pathSegRay, float& t_min, Geom *s_geoms, + int &hit_geom_index, glm::vec3 &normal) { + for (int i = 0; i < geoms_size; i++) { + float t; + glm::vec3 tmp_intersect; + glm::vec3 tmp_normal; + Geom & geom = s_geoms[i]; + + if (geom.type == CUBE) { + t = boxIntersectionTest(geom, /*pathSegment.ray*/ pathSegRay, tmp_intersect, tmp_normal/*, + outside*/); + } else if (geom.type == SPHERE) { + t = sphereIntersectionTest(geom, /*pathSegment.ray*/ pathSegRay, tmp_intersect, tmp_normal/*, + outside*/); + } + // TODO: add more intersection tests here... triangle? metaball? CSG? + + // Compute the minimum t from the intersection tests to determine what + // scene geometry object was hit first. + if (t > 0.0f && t_min > t) { + t_min = t; + hit_geom_index = i; + //intersect_point = tmp_intersect; + normal = tmp_normal; + } + } } // TODO: @@ -147,66 +233,40 @@ __global__ void generateRayFromCamera(Camera cam, int iter, int traceDepth, Path // Generating new rays is handled in your shader(s). // Feel free to modify the code below. __global__ void computeIntersections( - int depth - , int num_paths - , PathSegment * pathSegments - , Geom * geoms - , int geoms_size - , ShadeableIntersection * intersections - ) -{ + int depth, int num_paths, PathSegment * pathSegments, Geom * geoms, + int geoms_size, ShadeableIntersection * intersections) { + extern __shared__ Geom s_geoms[]; + int path_index = blockIdx.x * blockDim.x + threadIdx.x; - if (path_index < num_paths) - { - PathSegment pathSegment = pathSegments[path_index]; + if (threadIdx.x < geoms_size) { + s_geoms[threadIdx.x] = geoms[threadIdx.x]; + } + + __syncthreads(); + + if (path_index < num_paths) { + //PathSegment pathSegment = pathSegments[path_index]; + Ray pathSegRay = pathSegments[path_index].ray; - float t; - glm::vec3 intersect_point; + //glm::vec3 intersect_point; glm::vec3 normal; float t_min = FLT_MAX; int hit_geom_index = -1; - bool outside = true; + //bool outside = true; - glm::vec3 tmp_intersect; - glm::vec3 tmp_normal; // naive parse through global geoms - for (int i = 0; i < geoms_size; i++) - { - Geom & geom = geoms[i]; - - if (geom.type == CUBE) - { - t = boxIntersectionTest(geom, pathSegment.ray, tmp_intersect, tmp_normal, outside); - } - else if (geom.type == SPHERE) - { - t = sphereIntersectionTest(geom, pathSegment.ray, tmp_intersect, tmp_normal, outside); - } - // TODO: add more intersection tests here... triangle? metaball? CSG? - - // Compute the minimum t from the intersection tests to determine what - // scene geometry object was hit first. - if (t > 0.0f && t_min > t) - { - t_min = t; - hit_geom_index = i; - intersect_point = tmp_intersect; - normal = tmp_normal; - } - } + FindNearestIntersection(geoms_size, pathSegRay, t_min, s_geoms, + hit_geom_index, normal); - if (hit_geom_index == -1) - { + if (hit_geom_index == -1) { intersections[path_index].t = -1.0f; - } - else - { + } else { //The ray hits something intersections[path_index].t = t_min; - intersections[path_index].materialId = geoms[hit_geom_index].materialid; + intersections[path_index].materialId = s_geoms[hit_geom_index].materialid; intersections[path_index].surfaceNormal = normal; } } @@ -222,16 +282,11 @@ __global__ void computeIntersections( // Your shaders should handle that - this can allow techniques such as // bump mapping. __global__ void shadeFakeMaterial ( - int iter - , int num_paths - , ShadeableIntersection * shadeableIntersections - , PathSegment * pathSegments - , Material * materials - ) -{ + int iter, int num_paths, ShadeableIntersection * shadeableIntersections, + PathSegment * pathSegments, Material * materials) { + int idx = blockIdx.x * blockDim.x + threadIdx.x; - if (idx < num_paths) - { + if (idx < num_paths && pathSegments[idx].remainingBounces > 0) { ShadeableIntersection intersection = shadeableIntersections[idx]; if (intersection.t > 0.0f) { // if the intersection exists... // Set up the RNG @@ -240,20 +295,32 @@ __global__ void shadeFakeMaterial ( thrust::default_random_engine rng = makeSeededRandomEngine(iter, idx, 0); thrust::uniform_real_distribution u01(0, 1); - Material material = materials[intersection.materialId]; + Material &material = materials[intersection.materialId]; glm::vec3 materialColor = material.color; // If the material indicates that the object was a light, "light" the ray + PathSegment &seg = pathSegments[idx]; if (material.emittance > 0.0f) { - pathSegments[idx].color *= (materialColor * material.emittance); + seg.color *= (materialColor * material.emittance); + seg.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 { - float lightTerm = glm::dot(intersection.surfaceNormal, glm::vec3(0.0f, 1.0f, 0.0f)); + /*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 + pathSegments[idx].color *= u01(rng); // apply some noise because why not*/ + //PathSegment &seg = pathSegments[idx]; + if (material.hasReflective) { + scatterRaySpecular(seg, + seg.ray.origin + seg.ray.direction*intersection.t, + intersection.surfaceNormal, material, rng); + } else { + scatterRayUniform(seg, + seg.ray.origin + seg.ray.direction*intersection.t, + intersection.surfaceNormal, material, rng); + } } // If there was no intersection, color the ray black. // Lots of renderers use 4 channel color, RGBA, where A = alpha, often @@ -265,6 +332,70 @@ __global__ void shadeFakeMaterial ( } } +__global__ void DirectLighting(int num_paths, + PathSegment * pathSegments, + int *light_indices, + int lights_size, + Geom *geoms, + int geoms_size, + Material *materials, + int iter) { + extern __shared__ Geom s_geoms[]; + + int path_index = blockIdx.x * blockDim.x + threadIdx.x; + + if (threadIdx.x < geoms_size) { + s_geoms[threadIdx.x] = geoms[threadIdx.x]; + } + __syncthreads(); + + if (path_index < num_paths && pathSegments[path_index].remainingBounces >= 0) { + //if (path_index < num_paths) { + thrust::default_random_engine rng = makeSeededRandomEngine(iter, path_index, 0); + thrust::uniform_real_distribution u01(0, 1); + int light_index = light_indices[int(u01(rng)*lights_size)]; + Geom &geom = s_geoms[light_index]; + glm::vec3 coord; + if (geom.type == CUBE) { + int dir = u01(rng) * 3; + float sign = u01(rng) > 0.5 ? 1 : -1; + glm::vec2 offset(u01(rng), u01(rng)); + switch (dir) { + case 0: + coord = glm::vec3(sign, offset.x, offset.y); + break; + case 1: + coord = glm::vec3(offset.x, sign, offset.y); + break; + case 2: + coord = glm::vec3(offset.x, offset.y, sign); + break; + } + } else if (geom.type == SPHERE) { + float t = u01(rng); + float u = 2 * PI*u01(rng); + float v = sqrt(t*(1 - t)); + coord = glm::vec3(2 * v*cos(u), 1 - 2 * t, 2 * v*sin(u)); + } + glm::vec3 world_coord = multiplyMV(geom.transform, glm::vec4(coord, 1.0f)); + int hit_geom_index; + { + float t_min; + glm::vec3 normal; + Ray r = pathSegments[path_index].ray; + r.direction = r.origin - world_coord; + + FindNearestIntersection(geoms_size, r, t_min, + s_geoms, hit_geom_index, normal); + } + if (hit_geom_index == light_index) { + pathSegments[path_index].color *= materials[geom.materialid].color * materials[geom.materialid].emittance; + } else { + pathSegments[path_index].color = glm::vec3(0.0f); + } + } +} + // Add the current iteration's output to the overall image __global__ void finalGather(int nPaths, glm::vec3 * image, PathSegment * iterationPaths) { @@ -277,6 +408,32 @@ __global__ void finalGather(int nPaths, glm::vec3 * image, PathSegment * iterati } } +/*__global__ void is_terminated(int n, int *bools, ShadeableIntersection *inters) { + int index = (blockIdx.x * blockDim.x) + threadIdx.x; + if (index < n) { + bools[index] = inters[index].t < 0.0 ? 0 : 1; + } +}*/ +struct is_terminated { + __device__ bool operator()(const PathSegment &seg) { + return seg.color == glm::vec3(0.0f); + } +}; + +//struct order_intersections { +// __device__ bool operator()(const PtrTuple &a, const PtrTuple &b) { +// return thrust::get<1>(a).get()->materialId < +// thrust::get<1>(b).get()->materialId; +// } +//}; + +__global__ void extract_material(int n, int *indices, const ShadeableIntersection *inters) { + int index = threadIdx.x + blockIdx.x * blockDim.x; + if (index < n) { + indices[index] = inters->materialId; + } +} + /** * Wrapper for the __global__ call that sets up the kernel calls and does a ton * of memory management @@ -332,6 +489,8 @@ void pathtrace(uchar4 *pbo, int frame, int iter) { int depth = 0; PathSegment* dev_path_end = dev_paths + pixelcount; int num_paths = dev_path_end - dev_paths; + //num_paths = num_paths / 2-400; + int num_curr_paths = num_paths; // --- PathSegment Tracing Stage --- // Shoot ray into scene, bounce between objects, push shading chunks @@ -339,46 +498,94 @@ void pathtrace(uchar4 *pbo, int frame, int iter) { bool iterationComplete = false; while (!iterationComplete) { - // clean shading chunks - cudaMemset(dev_intersections, 0, pixelcount * sizeof(ShadeableIntersection)); - - // tracing - dim3 numblocksPathSegmentTracing = (num_paths + blockSize1d - 1) / blockSize1d; - computeIntersections <<>> ( - depth - , num_paths - , dev_paths - , dev_geoms - , hst_scene->geoms.size() - , dev_intersections - ); - checkCUDAError("trace one bounce"); - cudaDeviceSynchronize(); - depth++; - - - // TODO: - // --- Shading Stage --- - // Shade path segments based on intersections and generate new rays by - // evaluating the BSDF. - // Start off with just a big kernel that handles all the different - // materials you have in the scenefile. - // TODO: compare between directly shading the path segments and shading - // path segments that have been reshuffled to be contiguous in memory. - - shadeFakeMaterial<<>> ( - iter, - num_paths, - dev_intersections, - dev_paths, - dev_materials - ); - iterationComplete = true; // TODO: should be based off stream compaction results. + // clean shading chunks + cudaMemset(dev_intersections, 0, + pixelcount * sizeof(ShadeableIntersection)); + + // tracing + dim3 numblocksPathSegmentTracing = + (num_curr_paths + blockSize1d - 1) / blockSize1d; +#if CACHE_FIRST_BOUNCE + if (!first_bounce_cached && depth == 0) { + computeIntersections<<geoms.size() + sizeof(Ray) * blockSize1d>>> ( + depth, num_curr_paths, dev_paths, dev_geoms, hst_scene->geoms.size(), + dev_intersections); + first_bounce_cached = true; + cudaMemcpy(dev_first_bounce, dev_intersections, sizeof(ShadeableIntersection) * pixelcount, cudaMemcpyDeviceToDevice); + } else if (depth != 0){ + computeIntersections<<geoms.size() + sizeof(Ray) * blockSize1d>>> ( + depth, num_curr_paths, dev_paths, dev_geoms, hst_scene->geoms.size(), + dev_intersections); + } else { + cudaMemcpy(dev_intersections, dev_first_bounce, sizeof(ShadeableIntersection) * pixelcount, cudaMemcpyDeviceToDevice); + } +#else + computeIntersections<<geoms.size() + sizeof(Ray) * blockSize1d>>> ( + depth, num_curr_paths, dev_paths, dev_geoms, hst_scene->geoms.size(), + dev_intersections); +#endif + checkCUDAError("trace one bounce"); + cudaDeviceSynchronize(); + depth++; + printf("depth: %d\n", 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. + + // SORT +#if SORT_MATERIALS + auto zipped = thrust::make_zip_iterator(thrust::make_tuple(dev_thrust_paths, dev_thrust_intersections)); + extract_material<<>> ( + num_curr_paths, dev_material_indices, dev_intersections); + thrust::sort_by_key(dev_thrust_material_indices, + dev_thrust_material_indices + num_curr_paths, + zipped); +#endif + + shadeFakeMaterial<<>> ( + iter, + num_curr_paths, + dev_intersections, + dev_paths, + dev_materials + ); + //PathSegment *seg = (PathSegment*)malloc(sizeof(PathSegment)*num_paths); + //cudaMemcpy(seg, dev_paths, sizeof(PathSegment)*num_paths, cudaMemcpyDeviceToHost); + //printf("remaining: %d\n", seg[0].remainingBounces); + + thrust::device_ptr thrust_paths(dev_paths); + thrust::device_ptr new_end = thrust::remove_if( + thrust_paths, + thrust_paths + num_curr_paths, + is_terminated()); + num_curr_paths = new_end - thrust_paths; + + //iterationComplete = num_curr_paths == 0 ? true : false; // TODO: should be based off stream compaction results. + + //PathSegment *seg2 = (PathSegment*)malloc(sizeof(PathSegment)*num_paths); + //cudaMemcpy(seg2, dev_paths, sizeof(PathSegment)*num_paths, cudaMemcpyDeviceToHost); + //printf("remaining: %d\n", seg2[0].remainingBounces); + //free(seg); + //free(seg2); + printf("paths: %d\n", num_curr_paths); + iterationComplete = depth >= traceDepth ? true : false; } + dim3 numblocksPathSegmentTracing = + (num_curr_paths + blockSize1d - 1) / blockSize1d; + DirectLighting<<geoms.size()>>>( + num_curr_paths, dev_paths, dev_light_indices, num_lights, dev_geoms, hst_scene->geoms.size(), + dev_materials, iter); // Assemble this iteration and apply it to the image dim3 numBlocksPixels = (pixelcount + blockSize1d - 1) / blockSize1d; - finalGather<<>>(num_paths, dev_image, dev_paths); + finalGather<<>>(num_curr_paths, dev_image, dev_paths); /////////////////////////////////////////////////////////////////////////// diff --git a/src/scene.cpp b/src/scene.cpp index cbae043..fc051d7 100644 --- a/src/scene.cpp +++ b/src/scene.cpp @@ -116,15 +116,21 @@ int Scene::loadCamera() { string line; utilityCore::safeGetline(fp_in, line); + camera.focalLength = 5.0f; + camera.lensDiameter = 0.0f; while (!line.empty() && fp_in.good()) { vector tokens = utilityCore::tokenizeString(line); if (strcmp(tokens[0].c_str(), "EYE") == 0) { - camera.position = glm::vec3(atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str())); + camera.position = glm::vec3(atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str())); } else if (strcmp(tokens[0].c_str(), "LOOKAT") == 0) { - camera.lookAt = glm::vec3(atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str())); + 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())); - } + 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(), "FOCALLENGTH") == 0) { + camera.focalLength = atof(tokens[1].c_str()); + } else if (strcmp(tokens[0].c_str(), "LENSRADIUS") == 0) { + camera.lensDiameter = atof(tokens[1].c_str()); + } utilityCore::safeGetline(fp_in, line); } diff --git a/src/sceneStructs.h b/src/sceneStructs.h index b38b820..afcde43 100644 --- a/src/sceneStructs.h +++ b/src/sceneStructs.h @@ -49,6 +49,8 @@ struct Camera { glm::vec3 right; glm::vec2 fov; glm::vec2 pixelLength; + float focalLength; + float lensDiameter; }; struct RenderState { diff --git a/stream_compaction/CMakeLists.txt b/stream_compaction/CMakeLists.txt deleted file mode 100644 index ac358c9..0000000 --- a/stream_compaction/CMakeLists.txt +++ /dev/null @@ -1,7 +0,0 @@ -set(SOURCE_FILES - ) - -cuda_add_library(stream_compaction - ${SOURCE_FILES} - OPTIONS -arch=sm_20 - )