diff --git a/CMakeLists.txt b/CMakeLists.txt index 414ab4d..db40597 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -68,7 +68,7 @@ if(${CMAKE_SYSTEM_NAME} MATCHES "Darwin") endif() include_directories(.) -#add_subdirectory(stream_compaction) # TODO: uncomment if using your stream compaction +add_subdirectory(stream_compaction) # TODO: uncomment if using your stream compaction add_subdirectory(src) cuda_add_executable(${CMAKE_PROJECT_NAME} @@ -78,7 +78,7 @@ cuda_add_executable(${CMAKE_PROJECT_NAME} target_link_libraries(${CMAKE_PROJECT_NAME} src - #stream_compaction # TODO: uncomment if using your stream compaction + stream_compaction # TODO: uncomment if using your stream compaction ${CORELIBS} ) diff --git a/README.md b/README.md index 110697c..bf72dbf 100644 --- a/README.md +++ b/README.md @@ -3,11 +3,257 @@ 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) +* Michael Willett +* Tested on: Windows 10, I5-4690k @ 3.50GHz 8.00GB, GTX 750-TI 2GB (Personal Computer) -### (TODO: Your README) +## Contents +1. [Introduction](#intro) +2. [Basic Path Tracing](#part1) +3. [Additional Features](#part2) +4. [Performance](#performance) +5. [Build Instructions](#appendix) -*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. + +## Introduction: Scene Rendering +While many topics in computer graphics involve fast and effecient methods for improving real-time rendering, +the methods employed are often clever hacks on how to simulate natural looking scenes. Practice has shown, however, +that in order to generate the most realistic looking images, the developer must simulate the physical effects of how +light interacts with a scene. In the real world, light rays will emit from a source and bounce around and through the +objects in the environment, interacting with the colors of each object, until it finally hits the observer's eye. +This project is aimed to implement a basic path tracer that replicates these physics, resulting in the following image: + +Final Path Tracer +
+**Final image after 10,000 iterations. Spheres show perfect specular white,**
+**high specular black, imperfect specular gray, and clear refraction.** +
+ + +## Section 1: Path Tracing Algorithm +In the real world, what we see is a combination of an uncountable number of photons being emitted from every light source. +For computational efficiency, we choose to simulate the light rays in the reverse order; that is, we start by casting +simulated rays from the camera into each pixel of the scene. In a perfect world, the simulation of each photon would bounce +indefinitely until it hits a light source, however, since we do have limitations in the hardware, so we restrict the total number +of bounces. In this project, we assume each ray is allowed up to eight total bounces, with an additional direct lighting bounce +if enabled. + +In the initial implementation, materials followed three different light bouncing patterns: emissive light source, perfect reflection, +and perfect diffuse. The image below shows eight bounces with all the walls being perfectly diffuse (a ray can bounce in any random direction), +the sphere being perfectly reflective (the ray is perfectly reflected around the surface normal), and a single light source. The final image +generated after a total 5000 iterations was performed to addiquitely sample enough simulated photons to reflect a "real" image. The basic +algorithm runs as follows: + +1. Propogate new rays into the scene based off of last intersection. +2. Determine if the ray insersects with any objects, and if so select the closest intersection. +3. Update the color of the ray based on the material, and terminate the ray if the material is emissive. +4. Repeat steps 1-3 until the desired number of bounces has been reached. + +Basic Path Tracer +
**Basic depth-8 path tracing with only perfect diffuse and perfect specular materials**
+ + +## Section 2: Additional Features + +### Stream Compaction +The first attempted optimization was to use stream compaction to prevent the intersection and shader kernels from operating on unnecessary rays. +Implementations were done both using the THRUST library as well a custom compaction kernel, however, significant bottlenecks were observered as +can be seen in the table [here](#performance). The leading cause of the performance loss is due to the substatial amounts of memory reads/writes, +as simply removing the reording of the path array (but keeping index compaction) shows significant time improvements. It is possible that scenes +with a higher likelihood of early ray termination might see some performance improvement due to fewer blocks, but not enough rays were exiting to +be useful for this application. Additional effort was put into using shared memory for the custom implementation, however, during development this +method was failing for large arrays, and was never fully bench marked. Residual code can be found in the stream_compaction folder for future work. + +### First Bounce Cache +One performance oriented modification that did yield noticeable improvements was simply to store the first ray projection from the camera and intersections. +Theoretically this should have resulted in close to a 12% improvement, however, viewing the details of the performance analysis shows that ray generation +overhead was minimal. This leaves the intersection test as the only room for improvement, and while it was the slowest running kernel in the code, the +shading kernel was a close second, so unfortunately runtime improvement was limited to 7%. + +### Sorting Rays by Material +A final hypothesis for improving run time was to sort all rays by the material of the closest intersection. The idea is that kernel blocks would be +running the same calculations for successive rays, so there would be better kernel saturation per block. Unfortunately like the stream compaction results, +the overhead of rearranging the rays in the middle of the function resulted in significant overhead that was not compensated by the improvement to the +individual shader kernels. + +### Imperfect Specular Reflection +Aside from runtime improvements, there were several modifications to the physics of the ray bounces to generate more realistic effects. The first was to include +imperfect specular reflection to perform some random dispersion of the light ray relative to the perfect reflection, but not full hemisphere sampling. The random +ray was calculated using two random angles: rotation around the reflected vector θ, and the angle between the perfect reflection and the random vector φ. +θ was generated as a uniform sampling from 0 to 2π to reflect radially in any direction. The calculation for φ is slightly more complicated as it +determins how close to the perfect reflection the ray is scattered. Specifically, φ = acos(x1/(n+1)) , where n is the specular exponent of +the material, and x is randomly generated between 0 and 1. This results in random sampling in the hemisphere for n=0, and samples increasingly close to perfect +reflection as n grows larger. In the image below, we can see the specular sphere from the basic implementation was modified to have an imperfect specular reflection +with an exponent of 5. + +Imperfect Reflection +
**Imperfect specular sphere with specular exponent = 5**
+ +### Light Refraction with Fresnel Effects +The second major physics improvement was to add light refraction for entering and exiting tranlucent materials. Most students who have taken basic a basic +physics class know that light will bend slightly when entering a new material based off of its index of refraction according to +[Snell's Law](https://en.wikipedia.org/wiki/Snell%27s_law). Here, we can see a simulation of a simulation of a blue tinted window with internel reflection +on the edges: + +Refraction no-fresnel +
**Light refraction with internal reflection**
+ +Taking refraction one step further, we know that transparent surfaces will also show some minor reflection of the scene in addition to the transparent effects. +This is known as a Fresnel effect, and in this implementation it can be appoximated simply by randomly sampling a ray bounce as either a reflective or refractive +bounce relative to a likelihood calculated using [Schlick's Approximation](https://en.wikipedia.org/wiki/Schlick%27s_approximation). Below, you can see side by side +the simulation of a pool with and without Fresnel effects causing reflections of the scene on the pool water. + +
+Pool no-fresnel +Pool with fresnel +
+
**Pool of water with (right) and without (left) surface reflection**
+ +### Stochastic Sampled Anti-Aliasing +One noticeable effect of always casting the rays to the center of each pixel being rendered is that the first collisions always start from the same coordinates, +resulting in similar computiations for what is not normally a discrete parameter. The easiest way to solve this is to simply add a random offset of the exact +ray direction in the initial ray cast into the scene within the boundary of one pixel. In this case, we again used a uniform distribution within each individual +pixel, resulting in significant edge smoothing without the need to post process the image after the path tracer completes. Below shows the before and after + +anti-aliasing closeup +
**Anti-aliasing to noticeably smooth edge of black object**
+ + +### Direct Lighting +Another major limitation of the standard path tracing algorithm is that if a ray does not terminate at a light before the final depth is reached, it does not +accurately reflect the lighting conditions of the scene. The simplest approach to this is add a final step that terminates all rays by doing a final bounce to a +random point in a random emissive source in the scene. Without any additonal light bouncing, we can see the direct lighting alone can produce accurate shadows +and does a good job at providing crude lighting to the scene: + +Direct lighting effect +
**Direct lighting only**
+ +Once added as a final step of the standard depth 8 path tracer, we can now get significantly improved lighting for the scene in the same number of iterations +show below (left is with direct lighting). + +Depth-8 then DL +
**Final pass to add direct lighting (Left: direct lighting on. Right: reflected light only.)**
+ +Researching direct lighting online provided conflicting information on implementations. One of the ideas when computing direct lighting for a ray tracer instead +of a path tracer is to only compute direct lighting for the first bounce simillarly to the demo image, but then add that to the final path traced image. Unfortunately +this implementation resulted in significant blowout of the light sources and did not have the desired effect: + +Depth-8 plus DL +
**Single pass direct lighting added directly to indirect render**
+ + + +### Section 3: Performance Effects +All features discussed in Section 2 were analyzed relative to the baseline code for the simplest implementation. Unfortunately all but one attempt to improve performance by +better aligning rays to perform similar computations in the same blocks or to run fewer blocks ended up adding more overhead than they saved in resource. The only feature +that resulted in improvement in run speed was the first bounce cache, which resulted in a ~7% improvement. The shader kernel was still the second most time consuming process +after ray intersection calculations, so further improvements internal to that method could possibly improve GPU saturation. Unfortunately the method used for anti-aliasing +voids any time saved from caching first intersections, so if a smooth image is desired, we can not take advantage of that time save. + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + +
FeatureTime (ms)Change (%)
Default57.9n/a
Thrust Compaction119.6+106.6
Custom Compaction122.2+111.1
Material Sorting218.0+276.5
Cache First Bounce53.8-7.1
Anti-aliasing57.90.0
Direct Lighting61.0+5.4
+
*Tests measured using std::chrono in milliseconds*
+ + + +## Appendix: Build Instructions +* `src/` contains the source code. + +**CMake note:** Do not change any build settings or add any files to your +project directly (in Visual Studio, Nsight, etc.) Instead, edit the +`src/CMakeLists.txt` file. Any files you add must be added here. If you edit it, +just rebuild your VS/Nsight project to make it update itself. + +**If you experience linker errors on build related to the compute capability during thrust calls, edit the project to include the CUDA +library 'cudadevrt.lib'** + +#### Windows + +1. In Git Bash, navigate to your cloned project directory. +2. Create a `build` directory: `mkdir build` + * (This "out-of-source" build makes it easy to delete the `build` directory + and try again if something goes wrong with the configuration.) +3. Navigate into that directory: `cd build` +4. Open the CMake GUI to configure the project: + * `cmake-gui ..` or `"C:\Program Files (x86)\cmake\bin\cmake-gui.exe" ..` + * Don't forget the `..` part! + * Make sure that the "Source" directory is like + `.../Project3-Path-Tracer`. + * Click *Configure*. Select your version of Visual Studio, Win64. + (**NOTE:** you must use Win64, as we don't provide libraries for Win32.) + * If you see an error like `CUDA_SDK_ROOT_DIR-NOTFOUND`, + set `CUDA_SDK_ROOT_DIR` to your CUDA install path. This will be something + like: `C:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v7.5` + * Click *Generate*. +5. If generation was successful, there should now be a Visual Studio solution + (`.sln`) file in the `build` directory that you just created. Open this. + (from the command line: `explorer *.sln`) +6. Build. (Note that there are Debug and Release configuration options.) +7. Run. Make sure you run the `cis565_` target (not `ALL_BUILD`) by + right-clicking it and selecting "Set as StartUp Project". + * If you have switchable graphics (NVIDIA Optimus), you may need to force + your program to run with only the NVIDIA card. In NVIDIA Control Panel, + under "Manage 3D Settings," set "Multi-display/Mixed GPU acceleration" + to "Single display performance mode". + +#### OS X & Linux + +It is recommended that you use Nsight. + +1. Open Nsight. Set the workspace to the one *containing* your cloned repo. +2. *File->Import...->General->Existing Projects Into Workspace*. + * Select the Project 0 repository as the *root directory*. +3. Select the *cis565-* project in the Project Explorer. From the *Project* + menu, select *Build All*. + * For later use, note that you can select various Debug and Release build + configurations under *Project->Build Configurations->Set Active...*. +4. If you see an error like `CUDA_SDK_ROOT_DIR-NOTFOUND`: + * In a terminal, navigate to the build directory, then run: `cmake-gui ..` + * Set `CUDA_SDK_ROOT_DIR` to your CUDA install path. + This will be something like: `/usr/local/cuda` + * Click *Configure*, then *Generate*. +5. Right click and *Refresh* the project. +6. From the *Run* menu, *Run*. Select "Local C/C++ Application" and the + `cis565_` binary. diff --git a/img/AA_compare.PNG b/img/AA_compare.PNG new file mode 100644 index 0000000..130cd74 Binary files /dev/null and b/img/AA_compare.PNG differ diff --git a/img/all_effects.png b/img/all_effects.png new file mode 100644 index 0000000..0ca3b5e Binary files /dev/null and b/img/all_effects.png differ diff --git a/img/all_on.png b/img/all_on.png new file mode 100644 index 0000000..0ffac25 Binary files /dev/null and b/img/all_on.png differ diff --git a/img/antialiasing.png b/img/antialiasing.png new file mode 100644 index 0000000..e81034b Binary files /dev/null and b/img/antialiasing.png differ diff --git a/img/bad_random_sampling.png b/img/bad_random_sampling.png new file mode 100644 index 0000000..31ebcc1 Binary files /dev/null and b/img/bad_random_sampling.png differ diff --git a/img/bad_rng.png b/img/bad_rng.png new file mode 100644 index 0000000..d8900b2 Binary files /dev/null and b/img/bad_rng.png differ diff --git a/img/depth-8_perfect_spec.png b/img/depth-8_perfect_spec.png new file mode 100644 index 0000000..a35f198 Binary files /dev/null and b/img/depth-8_perfect_spec.png differ diff --git a/img/direct.png b/img/direct.png new file mode 100644 index 0000000..6af99ad Binary files /dev/null and b/img/direct.png differ diff --git a/img/direct_lighting.png b/img/direct_lighting.png new file mode 100644 index 0000000..7dccf92 Binary files /dev/null and b/img/direct_lighting.png differ diff --git a/img/direct_lighting_grayscale.png b/img/direct_lighting_grayscale.png new file mode 100644 index 0000000..c409802 Binary files /dev/null and b/img/direct_lighting_grayscale.png differ diff --git a/img/fresnel_working.png b/img/fresnel_working.png new file mode 100644 index 0000000..660c78e Binary files /dev/null and b/img/fresnel_working.png differ diff --git a/img/imperfect_specular.png b/img/imperfect_specular.png new file mode 100644 index 0000000..88c8f41 Binary files /dev/null and b/img/imperfect_specular.png differ diff --git a/img/no_fresnel.png b/img/no_fresnel.png new file mode 100644 index 0000000..91c64bd Binary files /dev/null and b/img/no_fresnel.png differ diff --git a/img/single_bounce_direct_lighting.png b/img/single_bounce_direct_lighting.png new file mode 100644 index 0000000..290920c Binary files /dev/null and b/img/single_bounce_direct_lighting.png differ diff --git a/img/working_refraction.png b/img/working_refraction.png new file mode 100644 index 0000000..a8942b7 Binary files /dev/null and b/img/working_refraction.png differ diff --git a/scenes/all_demo.txt b/scenes/all_demo.txt new file mode 100644 index 0000000..a53e143 --- /dev/null +++ b/scenes/all_demo.txt @@ -0,0 +1,171 @@ +// 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 +SPECEX 0 +RGB .85 .35 .35 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Diffuse green +MATERIAL 3 +RGB .35 .85 .35 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Specular white +MATERIAL 4 +RGB .98 .98 .98 +SPECEX 5 +SPECRGB .98 .98 .98 +REFL 1 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Specular black +MATERIAL 5 +RGB 0 0 0 +SPECEX 100 +SPECRGB .3 .3 .3 +REFL 1 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Refractive blue +MATERIAL 6 +RGB 1 1 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 1 +REFRIOR 1.6 +EMITTANCE 0 + +// Mirror +MATERIAL 7 +RGB 0 0 0 +SPECEX 1000 +SPECRGB 1 1 1 +REFL 1 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Camera +CAMERA +RES 800 800 +FOVY 45 +ITERATIONS 10000 +DEPTH 8 +FILE cornell2 +EYE 0.0 5 10.5 +LOOKAT 0 5 0 +UP 0 1 0 + + +// Ceiling light +OBJECT 0 +cube +material 0 +TRANS 0 10 0 +ROTAT 0 0 0 +SCALE 3 .3 6 + +// Floor +OBJECT 1 +cube +material 1 +TRANS 0 0 5 +ROTAT 0 0 0 +SCALE 10 .01 20 + +// Ceiling +OBJECT 2 +cube +material 1 +TRANS 0 10 5 +ROTAT 0 0 90 +SCALE .01 10 20 + +// 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 5 +ROTAT 0 0 0 +SCALE .01 10 20 + +// Right wall +OBJECT 5 +cube +material 3 +TRANS 5 5 5 +ROTAT 0 0 0 +SCALE .01 10 20 + +// Sphere +OBJECT 6 +sphere +material 4 +TRANS -2 2 0 +ROTAT 0 0 0 +SCALE 3 3 3 + +// Sphere +OBJECT 7 +sphere +material 6 +TRANS 2 2 0 +ROTAT 0 0 0 +SCALE 3 3 3 + +// Water +OBJECT 8 +sphere +material 5 +TRANS 2 6 -2 +ROTAT 0 0 0 +SCALE 3 3 3 + +// mirror sphere +OBJECT 9 +sphere +material 7 +TRANS -2 6 -2 +ROTAT 0 0 0 +SCALE 3 3 3 \ No newline at end of file diff --git a/scenes/dl.txt b/scenes/dl.txt new file mode 100644 index 0000000..5586ee6 --- /dev/null +++ b/scenes/dl.txt @@ -0,0 +1,83 @@ +// Emissive material (light) +MATERIAL 0 +RGB 1 1 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 3 + +// Diffuse white +MATERIAL 1 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Diffuse green +MATERIAL 2 +RGB .35 .85 .35 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Diffuse Gray +MATERIAL 3 +RGB .5 .5 .5 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Camera +CAMERA +RES 800 800 +FOVY 45 +ITERATIONS 5000 +DEPTH 8 +FILE direct_lighting +EYE 0.0 5 10.5 +LOOKAT 0 5 0 +UP 0 1 0 + + +// light +OBJECT 0 +sphere +material 0 +TRANS 10 10 0 +ROTAT 0 0 0 +SCALE 1 1 1 + +// Floor +OBJECT 1 +cube +material 3 +TRANS 0 0 0 +ROTAT 0 0 0 +SCALE 30 .01 30 + +// Left wall +OBJECT 2 +cube +material 2 +TRANS -5 5 0 +ROTAT 0 0 0 +SCALE .01 10 30 + +// Sphere +OBJECT 3 +sphere +material 1 +TRANS 0 4 0 +ROTAT 0 0 0 +SCALE 3 3 3 diff --git a/scenes/pool.txt b/scenes/pool.txt new file mode 100644 index 0000000..c2769a7 --- /dev/null +++ b/scenes/pool.txt @@ -0,0 +1,162 @@ +// 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 +SPECEX 0 +RGB .85 .35 .35 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Diffuse green +MATERIAL 3 +RGB .35 .85 .35 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Specular white +MATERIAL 4 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB .98 .98 .98 +REFL 1 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Specular black +MATERIAL 5 +RGB 0 0 0 +SPECEX 100 +SPECRGB .3 .3 .3 +REFL 1 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Refractive blue +MATERIAL 6 +RGB .85 .85 1.0 +SPECEX 100 +SPECRGB .9 .9 1.0 +REFL 0 +REFR 1 +REFRIOR 1.3 +EMITTANCE 0 + +// Camera +CAMERA +RES 800 800 +FOVY 45 +ITERATIONS 5000 +DEPTH 8 +FILE cornell2 +EYE 0.0 5 10.5 +LOOKAT 0 5 0 +UP 0 1 0 + + +// Ceiling light +OBJECT 0 +cube +material 0 +TRANS 0 10 0 +ROTAT 0 0 0 +SCALE 3 .3 3 + +// Floor +OBJECT 1 +cube +material 1 +TRANS 0 0 5 +ROTAT 0 0 0 +SCALE 10 .01 20 + +// Ceiling +OBJECT 2 +cube +material 1 +TRANS 0 10 5 +ROTAT 0 0 90 +SCALE .01 10 20 + +// 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 5 +ROTAT 0 0 0 +SCALE .01 10 20 + +// Right wall +OBJECT 5 +cube +material 3 +TRANS 5 5 5 +ROTAT 0 0 0 +SCALE .01 10 20 + +// Sphere +OBJECT 6 +sphere +material 4 +TRANS -1 4 -1 +ROTAT 0 0 0 +SCALE 3 3 3 + +// Sphere +OBJECT 7 +sphere +material 5 +TRANS 2 2 4 +ROTAT 0 0 0 +SCALE 2 2 2 + +// Water +OBJECT 8 +cube +material 6 +TRANS 0 1 5 +ROTAT 0 0 0 +SCALE 10 2 20 + + +// Front wall +OBJECT 9 +cube +material 2 +TRANS 0 5 15 +ROTAT 0 90 0 +SCALE .01 10 10 \ No newline at end of file diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index a1cb3fb..aaf8562 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_50 ) diff --git a/src/interactions.h b/src/interactions.h index d8107fb..b0ccb17 100644 --- a/src/interactions.h +++ b/src/interactions.h @@ -2,6 +2,8 @@ #include "intersections.h" +#define FRESNEL_EFFECTS 1 + // CHECKITOUT /** * Computes a cosine-weighted random direction in a hemisphere. @@ -41,6 +43,49 @@ glm::vec3 calculateRandomDirectionInHemisphere( + sin(around) * over * perpendicularDirection2; } + +/** +* Computes a random direction for imperfect specular reflection +*/ +__host__ __device__ +glm::vec3 calculateRandomDirectionReflective(glm::vec3 normal, thrust::default_random_engine &rng, const Material &m) +{ + thrust::uniform_real_distribution u01(0, 1); + + float n = m.specular.exponent; + float theta_s = acos(powf(u01(rng),1/(n+1))); + + float up = cos(theta_s); + float over = sin(theta_s); + float around = u01(rng) * TWO_PI; + + // Find a direction that is not the normal based off of whether or not the + // normal's components are all equal to sqrt(1/3) or whether or not at + // least one component is less than sqrt(1/3). Learned this trick from + // Peter Kutz. + + glm::vec3 directionNotNormal; + if (abs(normal.x) < SQRT_OF_ONE_THIRD) { + directionNotNormal = glm::vec3(1, 0, 0); + } + else if (abs(normal.y) < SQRT_OF_ONE_THIRD) { + directionNotNormal = glm::vec3(0, 1, 0); + } + else { + directionNotNormal = glm::vec3(0, 0, 1); + } + + // Use not-normal direction to generate two perpendicular directions + glm::vec3 perpendicularDirection1 = + glm::normalize(glm::cross(normal, directionNotNormal)); + glm::vec3 perpendicularDirection2 = + glm::normalize(glm::cross(normal, perpendicularDirection1)); + + return up * normal + + cos(around) * over * perpendicularDirection1 + + sin(around) * over * perpendicularDirection2; +} + /** * Scatter a ray with some probabilities according to the material properties. * For example, a diffuse surface scatters in a cosine-weighted hemisphere. @@ -70,11 +115,76 @@ __host__ __device__ void scatterRay( Ray &ray, glm::vec3 &color, - glm::vec3 intersect, - glm::vec3 normal, + const ShadeableIntersection intersect, 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 &l = ray.direction; + const glm::vec3 &n = intersect.surfaceNormal; + + + if (m.hasReflective > 0.0f) { + ray.origin = getPointOnRay(ray, intersect.t) + .01f * n; + ray.direction = calculateRandomDirectionReflective(glm::reflect(l, n), rng, m); + ray.indexOfRefraction = 1.0f; + color *= m.specular.color; + } + else if (m.hasRefractive > 0.0f) { + + float eta = (intersect.outside) ? 1.0f / m.indexOfRefraction : m.indexOfRefraction; + float k = 1.0f - eta * eta * (1.0 - glm::dot(-n, l) * glm::dot(-n, l)); + + if (k < 0) { //internal reflection + ray.origin = getPointOnRay(ray, intersect.t) + .01f * n; + ray.indexOfRefraction = m.indexOfRefraction; + ray.direction = glm::reflect(l, n); + color *= m.color; + } + else { + + if (intersect.outside) { + #if FRESNEL_EFFECTS + // Schlicks + float n1 = m.indexOfRefraction; + float r0 = (1.0f - n1) * (1.0f - n1) / ((1.0f + n1) * (1.0f + n1)); + float r = r0 + (1 - r0) * powf((1 - glm::dot(-n, l)), 5.0f); + assert(r <= 1.0f); + thrust::uniform_real_distribution u01(0, 1); + + if (u01(rng) < r) { // reflect + ray.origin = getPointOnRay(ray, intersect.t) + .01f * n; + ray.direction = glm::reflect(l, n); + ray.indexOfRefraction = 1.0f; + + color *= m.specular.color; + } + else + #endif + { // refract + ray.origin = getPointOnRay(ray, intersect.t) + .01f * -n; + ray.indexOfRefraction = m.indexOfRefraction; + ray.direction = glm::refract(l, n, eta); + color *= m.color; + } + + } + else { + ray.origin = getPointOnRay(ray, intersect.t) + .01f * -n; + ray.indexOfRefraction = m.indexOfRefraction; + ray.direction = glm::refract(l, n, eta); + color *= m.color; + } + } + + } + else { // assumes diffuse + ray.origin = getPointOnRay(ray, intersect.t) + .01f * n; + ray.direction = calculateRandomDirectionInHemisphere(n, rng); + ray.indexOfRefraction = 1.0f; + color *= m.color; + } + } diff --git a/src/pathtrace.cu b/src/pathtrace.cu index 94ccc42..2206bc7 100644 --- a/src/pathtrace.cu +++ b/src/pathtrace.cu @@ -1,10 +1,15 @@ #include #include #include +#include #include #include #include +#include +#include +#include +#include "stream_compaction\efficient.h" #include "sceneStructs.h" #include "scene.h" #include "glm/glm.hpp" @@ -14,6 +19,15 @@ #include "intersections.h" #include "interactions.h" + +#define SORT_BY_MATERIAL 0 +#define CACHE_FIRST_BOUNCE 0 +#define ANTIALIASING 1 +#define CUSTOM_COMPACT 0 +#define THRUST_COMPACT 0 +#define DIRECT_LIGHTING 1 + +#define BLOCK_SIZE 128 #define ERRORCHECK 1 #define FILENAME (strrchr(__FILE__, '/') ? strrchr(__FILE__, '/') + 1 : __FILE__) @@ -75,6 +89,12 @@ static PathSegment * dev_paths = NULL; static ShadeableIntersection * dev_intersections = NULL; // TODO: static variables for device memory, any extra info you need, etc // ... +static PathSegment * dev_first_paths = NULL; +static ShadeableIntersection * dev_first_intersection = NULL; +static Geom * dev_lights = NULL; +static int * dev_active = NULL; +static int * dev_inactive = NULL; +static int * dev_compacted = NULL; void pathtraceInit(Scene *scene) { hst_scene = scene; @@ -96,6 +116,21 @@ void pathtraceInit(Scene *scene) { cudaMemset(dev_intersections, 0, pixelcount * sizeof(ShadeableIntersection)); // TODO: initialize any extra device memeory you need + cudaMalloc(&dev_first_paths, pixelcount * sizeof(PathSegment)); + + cudaMalloc(&dev_first_intersection, pixelcount * sizeof(ShadeableIntersection)); + cudaMemset(dev_first_intersection, 0, pixelcount * sizeof(ShadeableIntersection)); + + cudaMalloc(&dev_lights, scene->geoms.size() * sizeof(Geom)); + + cudaMalloc(&dev_active, pixelcount * sizeof(int)); + cudaMemset(dev_active, 0, pixelcount * sizeof(int)); + + cudaMalloc(&dev_inactive, pixelcount * sizeof(int)); + cudaMemset(dev_inactive, 0, pixelcount * sizeof(int)); + + cudaMalloc(&dev_compacted, pixelcount * sizeof(int)); + cudaMemset(dev_compacted, 0, pixelcount * sizeof(int)); checkCUDAError("pathtraceInit"); } @@ -107,6 +142,12 @@ void pathtraceFree() { cudaFree(dev_materials); cudaFree(dev_intersections); // TODO: clean up any extra device memory you created + cudaFree(dev_first_paths); + cudaFree(dev_first_intersection); + cudaFree(dev_lights); + cudaFree(dev_active); + cudaFree(dev_inactive); + cudaFree(dev_compacted); checkCUDAError("pathtraceFree"); } @@ -129,16 +170,29 @@ __global__ void generateRayFromCamera(Camera cam, int iter, int traceDepth, Path PathSegment & segment = pathSegments[index]; segment.ray.origin = cam.position; - segment.color = glm::vec3(1.0f, 1.0f, 1.0f); + segment.color = glm::vec3(1.0f, 1.0f, 1.0f); + +#if ANTIALIASING + thrust::default_random_engine rng = makeSeededRandomEngine(index, iter, 0); + thrust::uniform_real_distribution u01(0, 1); + + glm::vec3 jitter((u01(rng) - 0.5f) * cam.pixelLength.x, (u01(rng) - 0.5f) * cam.pixelLength.y, 0); - // 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) + + jitter + ); +#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 segment.pixelIndex = index; segment.remainingBounces = traceDepth; + segment.ray.indexOfRefraction = 1.0f; } } @@ -152,10 +206,11 @@ __global__ void pathTraceOneBounce( ) { int path_index = blockIdx.x * blockDim.x + threadIdx.x; + PathSegment pathSegment; if (path_index < num_paths) { - PathSegment pathSegment = pathSegments[path_index]; + pathSegment = pathSegments[path_index]; float t; glm::vec3 intersect_point; @@ -163,34 +218,35 @@ __global__ void pathTraceOneBounce( float t_min = FLT_MAX; int hit_geom_index = -1; bool outside = true; - + + bool tmp_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); + t = boxIntersectionTest(geom, pathSegment.ray, tmp_intersect, tmp_normal, tmp_outside); } else if (geom.type == SPHERE) { - t = sphereIntersectionTest(geom, pathSegment.ray, tmp_intersect, tmp_normal, outside); + t = sphereIntersectionTest(geom, pathSegment.ray, tmp_intersect, tmp_normal, tmp_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. + // 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; + outside = tmp_outside; } } @@ -204,58 +260,153 @@ __global__ void pathTraceOneBounce( intersections[path_index].t = t_min; intersections[path_index].materialId = geoms[hit_geom_index].materialid; intersections[path_index].surfaceNormal = normal; + intersections[path_index].outside = outside; } } } -// LOOK: "fake" shader demonstrating what you might do with the info in -// a ShadeableIntersection, as well as how to use thrust's random number -// generator. Observe that since the thrust random number generator basically -// adds "noise" to the iteration, the image should start off noisy and get -// cleaner as more iterations are computed. -// -// 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 - , ShadeableIntersection * shadeableIntersections - , PathSegment * pathSegments - , Material * materials +__device__ PathSegment computeNewRay(int i, Material &material, PathSegment ¤tPath, ShadeableIntersection &intersection) +{ + // Set up the RNG + thrust::default_random_engine rng = makeSeededRandomEngine(currentPath.remainingBounces + currentPath.pixelIndex, i, intersection.materialId); + + // Create new path + PathSegment newPath = currentPath; + scatterRay(newPath.ray, newPath.color, intersection, material, rng); + newPath.remainingBounces = currentPath.remainingBounces - 1; + + return newPath; +} + + +__global__ void shadeMaterialSimple( + int iter, + int traceDepth, + 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 - 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. - 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) + { + PathSegment ¤tRay = pathSegments[idx]; + ShadeableIntersection intersection = shadeableIntersections[idx]; + + if (intersection.t > 0.0f) { // if the intersection exists... + 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) { + if (currentRay.remainingBounces > 0) { + currentRay.color *= (materialColor * material.emittance); + } + currentRay.remainingBounces = 0.0f; + } else if (currentRay.remainingBounces > 0) {; + currentRay = computeNewRay(iter, material, currentRay, intersection); + } else { + currentRay.color = glm::vec3(0.0f); // Bottomed out without hitting a light + } + + + } + else { + currentRay.color = glm::vec3(0.0f); + currentRay.remainingBounces = -1.0f; + } + } +} + +__global__ void directLighting( + int iter, + int num_paths, + ShadeableIntersection * shadeableIntersections, + PathSegment * pathSegments, + Material * materials, + Geom * geoms, + int geoms_size, + Geom * lights, + int lights_size + ) +{ + int path_index = blockIdx.x * blockDim.x + threadIdx.x; + PathSegment pathSegment; + + + if (path_index < num_paths) + { + PathSegment &pathSegment = pathSegments[path_index]; + + // choose random light + int l = -1; + thrust::default_random_engine rng = makeSeededRandomEngine(path_index, iter, l); + thrust::uniform_real_distribution u01(0, 1); + l = u01(rng)*lights_size; + + // generate a new ray to a random position on each light + Geom & geom = lights[l]; + glm::vec3 lightSample(u01(rng)*2.0f - 1.0f, u01(rng)*2.0f - 1.0f, u01(rng)*2.0f - 1.0f); + if (geom.type == CUBE) + { + lightSample = glm::normalize(lightSample) * geom.scale; + } + else if (geom.type == SPHERE) + { + lightSample *= geom.scale; + } + lightSample += geom.translation; + pathSegment.ray.direction = glm::normalize(lightSample - pathSegment.ray.origin); + pathSegment.ray.origin += 0.001f * pathSegment.ray.direction; + + // now check for intersections + float t; + glm::vec3 intersect_point; + glm::vec3 normal; + float t_min = FLT_MAX; + int hit_geom_index = -1; + bool outside = true; + + bool tmp_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, tmp_outside); + } + else if (geom.type == SPHERE) + { + t = sphereIntersectionTest(geom, pathSegment.ray, tmp_intersect, tmp_normal, tmp_outside); + } + + // Compute the minimum t from the intersection tests to determine what + // scene geometry object was hit first. + if (t > 0.0f && t_min > t) + { + t_min = t; + hit_geom_index = i; + intersect_point = tmp_intersect; + normal = tmp_normal; + outside = tmp_outside; + } + } + + if (hit_geom_index != -1) + { + //The ray hits something + Material &mat = materials[geoms[hit_geom_index].materialid]; + if (mat.emittance > 0.0f) + pathSegment.color *= (mat.color * mat.emittance); + } + } + } // Add the current iteration's output to the overall image @@ -270,6 +421,43 @@ __global__ void finalGather(int nPaths, glm::vec3 * image, PathSegment * iterati } } +__global__ void directGather(int nPaths, glm::vec3 * image, PathSegment * iterationPaths) +{ + int index = (blockIdx.x * blockDim.x) + threadIdx.x; + + if (index < nPaths) + { + PathSegment iterationPath = iterationPaths[index]; + image[iterationPath.pixelIndex] *= iterationPath.color; + } +} + +__global__ void getActiveRays(int nPaths, int * active, int *inactive, const PathSegment * iterationPaths) +{ + int index = (blockIdx.x * blockDim.x) + threadIdx.x; + + if (index < nPaths) + { + // NOTE: 1-index arrays so we don't compact element 0 + if (iterationPaths[index].remainingBounces > 0) + active[index] = index + 1; + else + inactive[index] = index + 1; + } +} + +__global__ void partitionRays(int nPathsPrev, int nPathsCompacted, int * active, int *inactive, PathSegment *out, const PathSegment *in) +{ + int index = (blockIdx.x * blockDim.x) + threadIdx.x; + + if (index < nPathsCompacted) + { + out[index] = in[active[index] - 1]; + } else if (index < nPathsPrev) { + out[index] = in[inactive[index - nPathsCompacted] - 1]; + } +} + /** * Wrapper for the __global__ call that sets up the kernel calls and does a ton * of memory management @@ -278,6 +466,7 @@ 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 materialcount = hst_scene->materials.size(); // 2D block for generating ray from camera const dim3 blockSize2d(8, 8); @@ -288,6 +477,12 @@ void pathtrace(uchar4 *pbo, int frame, int iter) { // 1D block for path tracing const int blockSize1d = 128; + int lightCount = 0; + for (int i = 0; i < hst_scene->geoms.size(); i++) { + if (hst_scene->materials[hst_scene->geoms[i].materialid].emittance > 0.0f) + cudaMemcpy(&dev_lights[lightCount++], &dev_geoms[i], sizeof(Geom), cudaMemcpyDeviceToDevice); + } + /////////////////////////////////////////////////////////////////////////// // Recap: @@ -319,59 +514,197 @@ void pathtrace(uchar4 *pbo, int frame, int iter) { // TODO: perform one iteration of path tracing - generateRayFromCamera <<>>(cam, iter, traceDepth, dev_paths); - checkCUDAError("generate camera ray"); - int depth = 0; PathSegment* dev_path_end = dev_paths + pixelcount; int num_paths = dev_path_end - dev_paths; + int num_paths_final = num_paths; + +#if CACHE_FIRST_BOUNCE && !ANTIALIASING + if (iter == 1) { + generateRayFromCamera << > >(cam, iter, traceDepth, dev_paths); + checkCUDAError("generate camera ray"); + cudaMemcpy(dev_first_paths, dev_paths, num_paths*sizeof(PathSegment), cudaMemcpyDeviceToDevice); + } + else { + cudaMemcpy(dev_paths, dev_first_paths, num_paths*sizeof(PathSegment), cudaMemcpyDeviceToDevice); + } +#else + generateRayFromCamera << > >(cam, iter, traceDepth, dev_paths); + checkCUDAError("generate camera ray"); +#endif + // --- PathSegment Tracing Stage --- // Shoot ray into scene, bounce between objects, push shading chunks - bool iterationComplete = false; + bool iterationComplete = false; + thrust::device_ptr path_ptr = thrust::device_pointer_cast(dev_paths); + thrust::device_ptr inter_ptr = thrust::device_pointer_cast(dev_intersections); + + // Shared Compaction Test + //int *dev_a, *dev_b; + //cudaMalloc((void**)&dev_a, 8 * sizeof(int)); + //cudaMalloc((void**)&dev_b, 8 * sizeof(int)); + //int tmp[8] = { 1, 0, 2, 4, 0, 5, 6, 7}; + //cudaMemcpy(dev_a, tmp, 8 * sizeof(int), cudaMemcpyHostToDevice); + // + //int rtn = StreamCompaction::Efficient::compact_dev(8, dev_b, dev_a); + //cudaMemcpy(tmp, dev_b, 8 * sizeof(int), cudaMemcpyDeviceToHost); + + //cudaFree(dev_a); + //cudaFree(dev_b); + while (!iterationComplete) { + // clean shading chunks + cudaMemset(dev_intersections, 0, pixelcount * sizeof(ShadeableIntersection)); - // clean shading chunks - cudaMemset(dev_intersections, 0, pixelcount * sizeof(ShadeableIntersection)); - - // tracing - dim3 numblocksPathSegmentTracing = (num_paths + blockSize1d - 1) / blockSize1d; - pathTraceOneBounce <<>> ( - 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. + // tracing + dim3 numblocksPathSegmentTracing = (num_paths + blockSize1d - 1) / blockSize1d; + #if CACHE_FIRST_BOUNCE && !ANTIALIASING + { + if (iter > 1 && depth == 0) { + cudaMemcpy(dev_intersections, dev_first_intersection, num_paths*sizeof(ShadeableIntersection), cudaMemcpyDeviceToDevice); + } else { + pathTraceOneBounce << > > ( + depth + , num_paths + , dev_paths + , dev_geoms + , hst_scene->geoms.size() + , dev_intersections + ); + + checkCUDAError("trace one bounce"); + cudaDeviceSynchronize(); + + if (iter == 1 && depth == 0) + cudaMemcpy(dev_first_intersection, dev_intersections, num_paths*sizeof(ShadeableIntersection), cudaMemcpyDeviceToDevice); + } + } + #else + { + pathTraceOneBounce << > > ( + depth + , num_paths + , dev_paths + , dev_geoms + , hst_scene->geoms.size() + , dev_intersections + ); + + checkCUDAError("trace one bounce"); + cudaDeviceSynchronize(); + } + #endif + + + 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. + + + //std::chrono::time_point start, sort; + //float sort_timer = 0; + //start = std::chrono::system_clock::now(); + + #if SORT_BY_MATERIAL + thrust::stable_sort_by_key(inter_ptr, inter_ptr + num_paths, path_ptr, MaterialCmp()); + #endif + + shadeMaterialSimple << > > ( + iter, + traceDepth, + num_paths, + dev_intersections, + dev_paths, + dev_materials + ); + +#if CUSTOM_COMPACT + // custom ray compaction + cudaMemset(dev_active, 0, pixelcount * sizeof(int)); + cudaMemset(dev_inactive, 0, pixelcount * sizeof(int)); + getActiveRays << > > (num_paths, dev_active, dev_inactive, dev_paths); + + // compact remaining ray indices + cudaDeviceProp prop; + cudaGetDeviceProperties(&prop, 0); + //printf(" threads: %i\n", prop.maxThreadsPerBlock); + +// int test_n = 512; +// int buf[512]; +// cudaMemcpy(buf, dev_active, test_n*sizeof(int), cudaMemcpyDeviceToHost); + int remainingPaths = StreamCompaction::Efficient::compact_dev(num_paths, dev_compacted, dev_active); + checkCUDAError("compact active rays"); + cudaMemcpy(dev_active, dev_compacted, num_paths*sizeof(int), cudaMemcpyDeviceToDevice); + //cudaMemcpy(buf, dev_active, test_n*sizeof(int), cudaMemcpyDeviceToHost); + + // compact completed ray indices + StreamCompaction::Efficient::compact_dev(num_paths, dev_compacted, dev_inactive); + cudaMemcpy(dev_inactive, dev_compacted, num_paths*sizeof(int), cudaMemcpyDeviceToDevice); + checkCUDAError("compact inactive rays"); + + + // sort final path objects + PathSegment *dev_out; + cudaMalloc((void**)&dev_out, pixelcount*sizeof(PathSegment)); + + partitionRays << > > (num_paths, remainingPaths, dev_active, dev_inactive, dev_out, dev_paths ); + cudaMemcpy(dev_paths, dev_out, num_paths*sizeof(PathSegment), cudaMemcpyDeviceToDevice); + checkCUDAError("partition rays"); + + + num_paths = remainingPaths; + cudaFree(dev_out); +#else + int remainingPaths = thrust::count_if(path_ptr, path_ptr + num_paths, TerminateRay()); + num_paths = (remainingPaths <= 0) ? 0 : pixelcount; +#endif + +#if THRUST_COMPACT + + // ray compaction with thrust + thrust::partition(path_ptr, path_ptr + num_paths, TerminateRay()); + num_paths_final = num_paths; + num_paths = thrust::count_if(path_ptr, path_ptr + num_paths, TerminateRay()); + +#endif + + iterationComplete = (num_paths <= 0); // TODO: should be based off stream compaction results. } - // Assemble this iteration and apply it to the image - dim3 numBlocksPixels = (pixelcount + blockSize1d - 1) / blockSize1d; - finalGather<<>>(num_paths, dev_image, dev_paths); + // Assemble this iteration and apply it to the image + dim3 numBlocksPixels = (pixelcount + blockSize1d - 1) / blockSize1d; + //finalGather << > >(pixelcount, dev_image, dev_paths); + +#if DIRECT_LIGHTING + // Direct Lighting + num_paths = pixelcount; + + directLighting << > > ( + iter, + num_paths_final, + dev_intersections, + dev_paths, + dev_materials, + dev_geoms, + hst_scene->geoms.size(), + dev_lights, + lightCount + ); + checkCUDAError("find direct lighting"); +#endif + + finalGather << > >(pixelcount, dev_image, dev_paths); + /////////////////////////////////////////////////////////////////////////// @@ -383,4 +716,4 @@ void pathtrace(uchar4 *pbo, int frame, int iter) { pixelcount * sizeof(glm::vec3), cudaMemcpyDeviceToHost); checkCUDAError("pathtrace"); -} +} \ No newline at end of file diff --git a/src/pathtrace.h b/src/pathtrace.h index 1241227..efab602 100644 --- a/src/pathtrace.h +++ b/src/pathtrace.h @@ -6,3 +6,17 @@ void pathtraceInit(Scene *scene); void pathtraceFree(); void pathtrace(uchar4 *pbo, int frame, int iteration); + +void scan_dev(int n, int *dev_data); + +inline int ilog2(int x) { + int lg = 0; + while (x >>= 1) { + ++lg; + } + return lg; +} + +inline int ilog2ceil(int x) { + return ilog2(x - 1) + 1; +} \ No newline at end of file diff --git a/src/preview.cpp b/src/preview.cpp index 4eb0bc1..0724977 100644 --- a/src/preview.cpp +++ b/src/preview.cpp @@ -1,5 +1,6 @@ #define _CRT_SECURE_NO_DEPRECATE #include +#include #include "main.h" #include "preview.h" @@ -169,9 +170,19 @@ bool init() { } void mainLoop() { + std::chrono::time_point start; + float dur; + float avg = 0.0f; + int i = 0; while (!glfwWindowShouldClose(window)) { glfwPollEvents(); - runCuda(); + + i++; + start = std::chrono::system_clock::now(); + runCuda(); + dur = (std::chrono::duration_cast (std::chrono::system_clock::now() - start)).count(); + avg += dur; + printf("Average time: %f\n", avg/i); string title = "CIS565 Path Tracer | " + utilityCore::convertIntToString(iteration) + " Iterations"; glfwSetWindowTitle(window, title.c_str()); @@ -185,6 +196,8 @@ void mainLoop() { glDrawElements(GL_TRIANGLES, 6, GL_UNSIGNED_SHORT, 0); glfwSwapBuffers(window); } + float time = (std::chrono::system_clock::now() - start).count(); + glfwDestroyWindow(window); glfwTerminate(); } diff --git a/src/sceneStructs.h b/src/sceneStructs.h index 13cc860..c4fa246 100644 --- a/src/sceneStructs.h +++ b/src/sceneStructs.h @@ -14,7 +14,8 @@ enum GeomType { struct Ray { glm::vec3 origin; - glm::vec3 direction; + glm::vec3 direction; + int indexOfRefraction; }; struct Geom { @@ -61,7 +62,7 @@ struct RenderState { struct PathSegment { Ray ray; - glm::vec3 color; + glm::vec3 color; int pixelIndex; int remainingBounces; }; @@ -78,4 +79,26 @@ struct ShadeableIntersection { float t; glm::vec3 surfaceNormal; int materialId; + bool outside; +}; + +struct TerminateRay { + __host__ __device__ + bool operator()(const PathSegment p) { + return p.remainingBounces > 0; + } }; + +struct MaterialCmp { + __host__ __device__ + bool operator()(const ShadeableIntersection p, const ShadeableIntersection q) { + return p.materialId < q.materialId; + } +}; + +struct IsLight { + __host__ __device__ + bool operator()(const Material g) { + return g.emittance > 0.0f; + } +}; \ No newline at end of file diff --git a/stream_compaction/CMakeLists.txt b/stream_compaction/CMakeLists.txt index ac358c9..74aa5bf 100644 --- a/stream_compaction/CMakeLists.txt +++ b/stream_compaction/CMakeLists.txt @@ -1,7 +1,11 @@ set(SOURCE_FILES + "common.h" + "common.cu" + "efficient.h" + "efficient.cu" ) cuda_add_library(stream_compaction ${SOURCE_FILES} - OPTIONS -arch=sm_20 + OPTIONS -arch=sm_50 ) diff --git a/stream_compaction/common.cu b/stream_compaction/common.cu new file mode 100644 index 0000000..ba1e3b6 --- /dev/null +++ b/stream_compaction/common.cu @@ -0,0 +1,38 @@ +#include "common.h" + +namespace StreamCompaction { +namespace Common { + +/** + * Maps an array to an array of 0s and 1s for stream compaction. Elements + * which map to 0 will be removed, and elements which map to 1 will be kept. + */ +__global__ void kernMapToBoolean(int n, int *bools, const int *idata) { + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= n) { + return; + } + + bools[index] = !(idata[index] == 0); +} + +/** + * Performs scatter on an array. That is, for each element in idata, + * if bools[idx] == 1, it copies idata[idx] to odata[indices[idx]]. + */ +__global__ void kernScatter(int n, int *odata, + const int *idata, const int *bools, const int *indices) { + + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= n) { + return; + } + + if (bools == NULL) odata[indices[index]] = idata[index]; + else if (bools[index] == 1) odata[indices[index]] = idata[index]; +} + + + +} +} diff --git a/stream_compaction/common.h b/stream_compaction/common.h new file mode 100644 index 0000000..fe47816 --- /dev/null +++ b/stream_compaction/common.h @@ -0,0 +1,31 @@ +#pragma once +#include + +#define FILENAME (strrchr(__FILE__, '/') ? strrchr(__FILE__, '/') + 1 : __FILE__) +#define checkCUDAError(msg) checkCUDAErrorFn(msg, FILENAME, __LINE__) + +#define BLOCK_SIZE 128 +#define MAX_BLOCKS 4096 + + +inline int ilog2(int x) { + int lg = 0; + while (x >>= 1) { + ++lg; + } + return lg; +} + +inline int ilog2ceil(int x) { + return ilog2(x - 1) + 1; +} + + +namespace StreamCompaction { +namespace Common { + __global__ void kernMapToBoolean(int n, int *bools, const int *idata); + + __global__ void kernScatter(int n, int *odata, + const int *idata, const int *bools, const int *indices); +} +} \ No newline at end of file diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu new file mode 100644 index 0000000..7e99942 --- /dev/null +++ b/stream_compaction/efficient.cu @@ -0,0 +1,206 @@ +#include +#include +#include +#include "common.h" +#include "efficient.h" + +#define SHARED_MEMORY 0 + +#define MAX_ARRAY_SIZE 1024 +#define NUM_BANKS 16 +#define LOG_NUM_BANKS 4 +#define CONFLICT_FREE_OFFSET(n) ((n) >> NUM_BANKS + (n) >> (2 * LOG_NUM_BANKS)) + +namespace StreamCompaction { +namespace Efficient { + + +__global__ void scanBlock(int n, int *odata, const int *idata) { + + extern __shared__ int temp[]; // allocated on invocation + + int blid = blockIdx.x * blockDim.x; + int thid = threadIdx.x; + + if (blid + thid >= n/2) { + return; + } + + int offset = 1; + + int ai = thid; + int bi = thid + (n / 2); + int bankOffsetA = CONFLICT_FREE_OFFSET(ai); + int bankOffsetB = CONFLICT_FREE_OFFSET(bi); + temp[ai + bankOffsetA] = idata[ai]; // load input into shared memory + temp[bi + bankOffsetB] = idata[bi]; + + for (int d = n >> 1; d > 0; d >>= 1) // build sum in place up the tree + { + __syncthreads(); + if (thid < d) + { + ai = offset*(2 * thid + 1) - 1; + bi = offset*(2 * thid + 2) - 1; + ai += CONFLICT_FREE_OFFSET(ai); + bi += CONFLICT_FREE_OFFSET(bi); + temp[bi] += temp[ai]; + } + offset *= 2; + } + + + if (thid == 0) { temp[n - 1 + CONFLICT_FREE_OFFSET(n - 1)] = 0; } // clear the last element + + for (int d = 1; d < n; d *= 2) // traverse down tree & build scan + { + offset >>= 1; + __syncthreads(); + if (thid < d) + { + ai = offset*(2 * thid + 1) - 1; + bi = offset*(2 * thid + 2) - 1; + ai += CONFLICT_FREE_OFFSET(ai); + bi += CONFLICT_FREE_OFFSET(bi); + + float t = temp[ai]; + temp[ai] = temp[bi]; + temp[bi] += t; + } + } + __syncthreads(); + + // write results to device memory + + odata[ai + blid] = temp[ai + bankOffsetA]; + odata[bi + blid] = temp[bi + bankOffsetB]; +} + +__global__ void scanMultipleBlocks(int n, int *odata, const int *idata) { +} + +__global__ void kernUpStep(int n, int d, int *data) { + + int s = 1 << (d + 1); + int index = (threadIdx.x + (blockIdx.x * blockDim.x)) *s; + + if (index >= n) { + return; + } + + data[index + s - 1] += data[index + s / 2 - 1]; +} + +__global__ void kernDownStep(int n, int d, int *data) { + + int s = 1 << (d + 1); + int index = (threadIdx.x + (blockIdx.x * blockDim.x)) * s; + + if (index >= n) { + return; + } + + + int t = data[index + s / 2 - 1]; + data[index + s / 2 - 1] = data[index + s - 1]; + data[index + s - 1] += t; +} + + +/** +* Performs prefix-sum (aka scan) on idata, storing the result into odata. +* For use with arrays intiialized on GPU already. +*/ +void scan_dev(int n, int *dev_in) { + + + +#if SHARED_MEMORY + // create device arrays to pad to power of 2 size array + //int pot = pow(2, ilog2ceil(n)); + int pot = 1 << ilog2ceil(n); + int *dev_data; + int host[512] = {0}; + + cudaMalloc((void**)&dev_data, pot*sizeof(int)); + cudaMemset(dev_data, 0, pot*sizeof(int)); + cudaMemcpy(dev_data, dev_in, n*sizeof(int), cudaMemcpyDeviceToDevice); + cudaMemcpy(host, dev_data, 256 * sizeof(int), cudaMemcpyDeviceToHost); + + dim3 fullBlocksPerGrid((n + MAX_ARRAY_SIZE - 1) / MAX_ARRAY_SIZE); + scanBlock << < fullBlocksPerGrid, MAX_ARRAY_SIZE, MAX_ARRAY_SIZE >> >(pot, dev_data, dev_in); + cudaMemcpy(host, dev_data, 512 * sizeof(int), cudaMemcpyDeviceToHost); + + cudaMemcpy(dev_in, dev_data, n*sizeof(int), cudaMemcpyDeviceToDevice); + cudaFree(dev_data); + +#else + dim3 fullBlocksPerGrid((n + BLOCK_SIZE - 1) / BLOCK_SIZE); + // create device arrays to pad to power of 2 size array + //int pot = pow(2, ilog2ceil(n)); + int pot = 1 << ilog2ceil(n); + int *dev_data; + cudaMalloc((void**)&dev_data, pot*sizeof(int)); + cudaMemset(dev_data, 0, pot*sizeof(int)); + cudaMemcpy(dev_data, dev_in, n*sizeof(int), cudaMemcpyDeviceToDevice); + + float d = 0; + for (d; d < ilog2ceil(pot); d++) { + int div = 1 << (int) d + 1; + fullBlocksPerGrid.x = ((pot / div + BLOCK_SIZE - 1) / BLOCK_SIZE); + kernUpStep << < fullBlocksPerGrid, BLOCK_SIZE >> >(pot, d, dev_data); + } + + cudaMemset(&dev_data[pot - 1], 0, sizeof(int)); + for (d = ilog2ceil(pot); d >= 0; d--) { + int div = 1 << (int) d + 1; + fullBlocksPerGrid.x = ((pot / div + BLOCK_SIZE - 1) / BLOCK_SIZE); + kernDownStep << < fullBlocksPerGrid, BLOCK_SIZE >> >(pot, d, dev_data); + } + + cudaMemcpy(dev_in, dev_data, n*sizeof(int), cudaMemcpyDeviceToDevice); + cudaFree(dev_data); +#endif + +} + + +int compact_dev(int n, int *dev_out, const int *dev_in) { + + // create device arrays + int *dev_indices; + int *dev_bools; + int rtn = -1; + + //int n_iter = 1 << 24; + + cudaMalloc((void**)&dev_indices, n*sizeof(int)); + cudaMalloc((void**)&dev_bools, n*sizeof(int)); + + dim3 fullBlocksPerGrid((n + BLOCK_SIZE - 1) / BLOCK_SIZE); + cudaMemset(dev_indices, 0, n); + cudaMemset(dev_bools, 0, n); + + StreamCompaction::Common::kernMapToBoolean << < fullBlocksPerGrid, BLOCK_SIZE >> >(n, dev_bools, dev_in); + + // scan without wasteful device-host-device write + cudaMemcpy(dev_indices, dev_bools, n*sizeof(int), cudaMemcpyDeviceToDevice); + scan_dev(n, dev_indices); + + // scatter + StreamCompaction::Common::kernScatter << < fullBlocksPerGrid, BLOCK_SIZE >> >(n, dev_out, dev_in, dev_bools, dev_indices); + cudaMemcpy(&rtn, &dev_indices[n - 1], sizeof(int), cudaMemcpyDeviceToHost); + + // Synch and throw away run time error that thinks we can't handle arrays over 2^16 + cudaDeviceSynchronize(); + cudaError_t err = cudaGetLastError(); + + cudaFree(dev_bools); + cudaFree(dev_indices); + + return rtn; +} + + +} +} diff --git a/stream_compaction/efficient.h b/stream_compaction/efficient.h new file mode 100644 index 0000000..df540f4 --- /dev/null +++ b/stream_compaction/efficient.h @@ -0,0 +1,8 @@ +#pragma once + +namespace StreamCompaction { +namespace Efficient { + void scan_dev(int n, int *dev_data); + int compact_dev(int n, int *dev_out, const int *dev_in); +} +} diff --git a/style.css b/style.css new file mode 100644 index 0000000..e9331a7 --- /dev/null +++ b/style.css @@ -0,0 +1,7 @@ +.center {margin: auto;display: block;} +.tg {border-collapse:collapse;border-spacing:0;border-color:#999;width:400px;margin: auto;display: block;} +.tg td{font-family:Arial, sans-serif;font-size:14px;padding:10px 5px;border-style:solid;border-width:0px;overflow:hidden;word-break:normal;border-color:#999;color:#444;background-color:#F7FDFA;} +.tg th{font-family:Arial, sans-serif;font-size:14px;font-weight:normal;padding:10px 5px;border-style:solid;border-width:0px;overflow:hidden;word-break:normal;border-color:#999;color:#fff;background-color:#26ADE4;} +.tg .tg-baqh{text-align:center;vertical-align:top} +.tg .tg-amwm{font-weight:bold;text-align:center;vertical-align:top} +.tg .tg-9ewa{color:#fe0000;text-align:center;vertical-align:top} \ No newline at end of file