diff --git a/README.md b/README.md index ae6088d..8118cdd 100644 --- a/README.md +++ b/README.md @@ -1,273 +1,89 @@ -CIS 565 Project3 : CUDA Pathtracer +CIS 565 project 03 : CUDA path tracer =================== -Fall 2014 +## INTRODUCTION -Due Wed, 10/8 (submit without penalty until Sun, 10/12) +This project is a CUDA-parallelized Monte Carlo path tracer implemented for CIS 565 during my fall 2014 semester at Penn. Given a scene file that defines a camera, materials, and geometry, my path tracer is capable of rendering images with full global illumination, including realistic diffuse surfaces, color bleeding, caustics, area lights, and soft shadows. Additionally, my path tracer supports Fresnel refractions for glass materials and texture mapping for both spheres and cubes. -## INTRODUCTION -In this project, you will implement a CUDA based pathtracer capable of -generating pathtraced rendered images extremely quickly. Building a pathtracer can be viewed as a generalization of building a raytracer, so for those of you who have taken 460/560, the basic concept should not be very new to you. For those of you that have not taken -CIS460/560, raytracing is a technique for generating images by tracing rays of -light through pixels in an image plane out into a scene and following the way -the rays of light bounce and interact with objects in the scene. More -information can be found here: -http://en.wikipedia.org/wiki/Ray_tracing_(graphics). Pathtracing is a generalization of this technique by considering more than just the contribution of direct lighting to a surface. - -Since in this class we are concerned with working in generating actual images -and less so with mundane tasks like file I/O, this project includes basecode -for loading a scene description file format, described below, and various other -things that generally make up the render "harness" that takes care of -everything up to the rendering itself. The core renderer is left for you to -implement. Finally, note that while this basecode is meant to serve as a -strong starting point for a CUDA pathtracer, you are not required to use this -basecode if you wish, and you may also change any part of the basecode -specification as you please, so long as the final rendered result is correct. - -## CONTENTS -The Project3 root directory contains the following subdirectories: - -* src/ contains the source code for the project. Both the Windows Visual Studio - solution and the OSX and Linux makefiles reference this folder for all - source; the base source code compiles on Linux, OSX and Windows without - modification. If you are building on OSX, be sure to uncomment lines 4 & 5 of - the CMakeLists.txt in order to make sure CMake builds against clang. -* data/scenes/ contains an example scene description file. -* renders/ contains an example render of the given example scene file. -* windows/ contains a Windows Visual Studio 2010 project and all dependencies - needed for building and running on Windows 7. If you would like to create a - Visual Studio 2012 or 2013 projects, there are static libraries that you can - use for GLFW that are in external/bin/GLFW (Visual Studio 2012 uses msvc110, - and Visual Studio 2013 uses msvc120) -* external/ contains all the header, static libraries and built binaries for - 3rd party libraries (i.e. glm, GLEW, GLFW) that we use for windowing and OpenGL - extensions - -## RUNNING THE CODE -The main function requires a scene description file (that is provided in data/scenes). -The main function reads in the scene file by an argument as such : -'scene=[sceneFileName]' - -If you are using Visual Studio, you can set this in the Debugging > Command Arguments section -in the Project properties. - -## REQUIREMENTS -In this project, you are given code for: - -* Loading, reading, and storing the scene scene description format -* Example functions that can run on both the CPU and GPU for generating random - numbers, spherical intersection testing, and surface point sampling on cubes -* A class for handling image operations and saving images -* Working code for CUDA-GL interop - -You will need to implement the following features: - -* Raycasting from a camera into a scene through a pixel grid -* Diffuse surfaces -* Perfect specular reflective surfaces -* Cube intersection testing -* Sphere surface point sampling -* Stream compaction optimization - -You are also required to implement at least 2 of the following features: - -* Texture mapping -* Bump mapping -* Depth of field -* Refraction, i.e. glass -* OBJ Mesh loading and rendering -* Interactive camera -* Motion blur -* Subsurface scattering - -The 'extra features' list is not comprehensive. If you have a particular feature -you would like to implement (e.g. acceleration structures, etc.) please contact us -first! - -For each 'extra feature' you must provide the following analysis : -* overview write up of the feature -* performance impact of the feature -* if you did something to accelerate the feature, why did you do what you did -* compare your GPU version to a CPU version of this feature (you do NOT need to - implement a CPU version) -* how can this feature be further optimized (again, not necessary to implement it, but - should give a roadmap of how to further optimize and why you believe this is the next - step) - -## BASE CODE TOUR -You will be working in three files: raytraceKernel.cu, intersections.h, and -interactions.h. Within these files, areas that you need to complete are marked -with a TODO comment. Areas that are useful to and serve as hints for optional -features are marked with TODO (Optional). Functions that are useful for -reference are marked with the comment LOOK. - -* raytraceKernel.cu contains the core raytracing CUDA kernel. You will need to - complete: - * cudaRaytraceCore() handles kernel launches and memory management; this - function already contains example code for launching kernels, - transferring geometry and cameras from the host to the device, and transferring - image buffers from the host to the device and back. You will have to complete - this function to support passing materials and lights to CUDA. - * raycastFromCameraKernel() is a function that you need to implement. This - function once correctly implemented should handle camera raycasting. - * raytraceRay() is the core raytracing CUDA kernel; all of your pathtracing - logic should be implemented in this CUDA kernel. raytraceRay() should - take in a camera, image buffer, geometry, materials, and lights, and should - trace a ray through the scene and write the resultant color to a pixel in the - image buffer. - -* intersections.h contains functions for geometry intersection testing and - point generation. You will need to complete: - * boxIntersectionTest(), which takes in a box and a ray and performs an - intersection test. This function should work in the same way as - sphereIntersectionTest(). - * getRandomPointOnSphere(), which takes in a sphere and returns a random - point on the surface of the sphere with an even probability distribution. - This function should work in the same way as getRandomPointOnCube(). You can - (although do not necessarily have to) use this to generate points on a sphere - to use a point lights, or can use this for area lighting. - -* interactions.h contains functions for ray-object interactions that define how - rays behave upon hitting materials and objects. You will need to complete: - * getRandomDirectionInSphere(), which generates a random direction in a - sphere with a uniform probability. This function works in a fashion - similar to that of calculateRandomDirectionInHemisphere(), which generates a - random cosine-weighted direction in a hemisphere. - * calculateBSDF(), which takes in an incoming ray, normal, material, and - other information, and returns an outgoing ray. You can either implement - this function for ray-surface interactions, or you can replace it with your own - function(s). - -You will also want to familiarize yourself with: - -* sceneStructs.h, which contains definitions for how geometry, materials, - lights, cameras, and animation frames are stored in the renderer. -* utilities.h, which serves as a kitchen-sink of useful functions - -## NOTES ON GLM -This project uses GLM, the GL Math library, for linear algebra. You need to -know two important points on how GLM is used in this project: - -* In this project, indices in GLM vectors (such as vec3, vec4), are accessed - via swizzling. So, instead of v[0], v.x is used, and instead of v[1], v.y is - used, and so on and so forth. -* GLM Matrix operations work fine on NVIDIA Fermi cards and later, but - pre-Fermi cards do not play nice with GLM matrices. As such, in this project, - GLM matrices are replaced with a custom matrix struct, called a cudaMat4, found - in cudaMat4.h. A custom function for multiplying glm::vec4s and cudaMat4s is - provided as multiplyMV() in intersections.h. - -## SCENE FORMAT -This project uses a custom scene description format. -Scene files are flat text files that describe all geometry, materials, -lights, cameras, render settings, and animation frames inside of the scene. -Items in the format are delimited by new lines, and comments can be added at -the end of each line preceded with a double-slash. - -Materials are defined in the following fashion: - -* MATERIAL (material ID) //material header -* RGB (float r) (float g) (float b) //diffuse color -* SPECX (float specx) //specular exponent -* SPECRGB (float r) (float g) (float b) //specular color -* REFL (bool refl) //reflectivity flag, 0 for - no, 1 for yes -* REFR (bool refr) //refractivity flag, 0 for - no, 1 for yes -* REFRIOR (float ior) //index of refraction - for Fresnel effects -* SCATTER (float scatter) //scatter flag, 0 for - no, 1 for yes -* ABSCOEFF (float r) (float b) (float g) //absorption - coefficient for scattering -* RSCTCOEFF (float rsctcoeff) //reduced scattering - coefficient -* EMITTANCE (float emittance) //the emittance of the - material. Anything >0 makes the material a light source. - -Cameras are defined in the following fashion: - -* CAMERA //camera header -* RES (float x) (float y) //resolution -* FOVY (float fovy) //vertical field of - view half-angle. the horizonal angle is calculated from this and the - reslution -* ITERATIONS (float interations) //how many - iterations to refine the image, only relevant for supersampled antialiasing, - depth of field, area lights, and other distributed raytracing applications -* FILE (string filename) //file to output - render to upon completion -* frame (frame number) //start of a frame -* EYE (float x) (float y) (float z) //camera's position in - worldspace -* VIEW (float x) (float y) (float z) //camera's view - direction -* UP (float x) (float y) (float z) //camera's up vector - -Objects are defined in the following fashion: -* OBJECT (object ID) //object header -* (cube OR sphere OR mesh) //type of object, can - be either "cube", "sphere", or "mesh". Note that cubes and spheres are unit - sized and centered at the origin. -* material (material ID) //material to - assign this object -* frame (frame number) //start of a frame -* TRANS (float transx) (float transy) (float transz) //translation -* ROTAT (float rotationx) (float rotationy) (float rotationz) //rotation -* SCALE (float scalex) (float scaley) (float scalez) //scale - -An example scene file setting up two frames inside of a Cornell Box can be -found in the scenes/ directory. - -For meshes, note that the base code will only read in .obj files. For more -information on the .obj specification see http://en.wikipedia.org/wiki/Wavefront_.obj_file. - -An example of a mesh object is as follows: - -OBJECT 0 -mesh tetra.obj -material 0 -frame 0 -TRANS 0 5 -5 -ROTAT 0 90 0 -SCALE .01 10 10 - -Check the Google group for some sample .obj files of varying complexity. - -## THIRD PARTY CODE POLICY -* Use of any third-party code must be approved by asking on our Google Group. - If it is approved, all students are welcome to use it. Generally, we approve - use of third-party code that is not a core part of the project. For example, - for the ray tracer, we would approve using a third-party library for loading - models, but would not approve copying and pasting a CUDA function for doing - refraction. -* Third-party code must be credited in README.md. -* Using third-party code without its approval, including using another - student's code, is an academic integrity violation, and will result in you - receiving an F for the semester. - -## SELF-GRADING -* On the submission date, email your grade, on a scale of 0 to 100, to Harmony, - harmoli+cis565@seas.upenn.com, with a one paragraph explanation. Be concise and - realistic. Recall that we reserve 30 points as a sanity check to adjust your - grade. Your actual grade will be (0.7 * your grade) + (0.3 * our grade). We - hope to only use this in extreme cases when your grade does not realistically - reflect your work - it is either too high or too low. In most cases, we plan - to give you the exact grade you suggest. -* Projects are not weighted evenly, e.g., Project 0 doesn't count as much as - the path tracer. We will determine the weighting at the end of the semester - based on the size of each project. - -## SUBMISSION -Please change the README to reflect the answers to the questions we have posed -above. Remember: -* this is a renderer, so include images that you've made! -* be sure to back your claims for optimization with numbers and comparisons -* if you reference any other material, please provide a link to it -* you wil not e graded on how fast your path tracer runs, but getting close to - real-time is always nice -* if you have a fast GPU renderer, it is good to show case this with a video to - show interactivity. If you do so, please include a link. - -Be sure to open a pull request and to send Harmony your grade and why you -believe this is the grade you should get. +## PARALLELIZATION SCHEME + +My path tracer is parallelized per ray rather than per pixel. At the start of each iteration, one ray is generated for each pixel in the image buffer and stored in a ray pool. At each trace depth (basically every time a ray intersects geometry), rays are checked to see if they should be retired from the ray pool. In my path tracer, rays are retired if they (A) do not intersect with any piece of geometry in the scene, or (B) intersect with a light source. A retired ray is removed from the ray pool and will not be considered during future kernel calls to the GPU. + +A per-ray parallelization scheme such as this prevents unwanted cases where some rays in a warp become inactive at a low trace depth while neighboring rays remain active until the max trace depth. In these circumstances, valuable GPU processing time is wasted on inactive rays that no longer contribute to the final rendered image result. + +## STREAM COMPACTION + +To support a per-ray parallelization scheme, stream compaction is used to cull retired rays from the ray pool. I use the thrust parallel algorithms library (https://code.google.com/p/thrust/) to perform my stream compaction. After all computations have been performed for rays at a certain trace depth, the ray pool is checked for retired rays (identified by a boolean member attached to each ray), and if a retired ray is found, I remove it from the ray pool with thrust's remove_if method. Then, during the next trace depth iteration, I send the current ray pool to the GPU which no longer contains any retired rays. + +## SUPPORTED GEOMETRY AND MATERIALS + +Currently, my path tracer supports sphere and cube geometry and ideal diffuse, perfectly specular, and glass materials. In the future, I plan to support arbitrary mesh objects and more complex BRDF models. + +## FRESNEL REFRACTION + +As can be seen in the image below, my path tracer supports glass materials. Transmission direction and probability of reflection and refraction are computed using the well-known Fresnel equations. An interesting implementation detail to note is that my parallelization scheme does not support recursion. This is due in part to my decision to never spawn new rays inside CUDA kernels which stems from my decision to initialize the ray pool with a fixed max size. In traditional path tracers and ray tracers, radiance at a point is often computed through recursion. This is especially important for refractive materials because when a ray intersects a refractive surface, often two rays are emitted from the intersection point--one reflective ray and one refractive ray. All these rays contribute to the final computed pixel color. + +In my path tracer, when rays intersect with a refractive surface, I use Fresnel's equations to determine the probability that a ray will either reflect or refract, but never both. In the future, I would like to explore methods in which I can retain my per-ray parallelization scheme while also dynamically spawning new rays when rays intersect refractive surfaces. + +In the image below, some green artifacts can be seen (most prominently on the left side where the green and white walls meet). I suspect these artifacts are caused by an integer overflow when seeding my random number generators because they only appear after a large number of iterations (~4000). However, I have not spent much time looking into this issue, so I cannot be sure of the cause. + +![alt tag](https://raw.githubusercontent.com/drerucha/Project3-Pathtracer/master/data/readme_pics/fresnel_refractions.jpg) + +## TEXTURE MAPPING + +I implemented texture mapping for spheres and cubes. When a piece of geometry with an applied texture is intersected, the intersection point is transformed to object-space where the uv-coordinates ([0, 1]) for the image are computed. These uv-coordinates are then multiplied by the texture dimensions to determine the desired pixel within the texture. It is this texture pixel's RGB information that is used when computing radiance instead of the geometry's material RGB. + +For me, the most difficult part of implementing texture mapping was getting the texture information onto the GPU. First, I tried passing in a list of data structures representing each texture, but I ran into problems allocating memory on the GPU since each texture could theoretically have different dimensions, and thus have different memory requirements. To use a list of data structures, I needed to define a fixed, uniform maximum texture resolution for each texture. This ensured that each texture structure had the same memory requirements, and made GPU memory allocation straightforward. However, defining a maximum texture resolution was too limiting. + +Instead, I opted to "flatten" all my textures into a single array of glm::vec3s representing RGB values at runtime. Additionally, I created an array of ints at runtime that stored the dimensions for each texture as well as the starting index within the array of RGB values for each texture. Through my work implementing texture mapping, I've learned that while CUDA supports multidimensional data and complex data types, it is most happy with flat data and primitive data types, and I was happy to honor its preferences. + +Currently, my textures do not support any kind of transparency, but I would like to utilize image alpha channels in the future. + +![alt tag](https://raw.githubusercontent.com/drerucha/Project3-Pathtracer/master/data/readme_pics/texture_mapping.jpg) + +![alt tag](https://raw.githubusercontent.com/drerucha/Project3-Pathtracer/master/data/readme_pics/texture_mapping_02.jpg) + +## JITTERED SUPERSAMPLED ANTI-ALIASING + +Supersampled anti-aliasing is a method to remove jagged edges that involves averaging multiple samples for every pixel where each sample (ray) originates from a different location from within a pixel (not just from the pixel's center). With a path tracer, since many rays (100+) must be processed for each pixel in order to create a physically plausible image, supersampled anti-aliasing comes for "free". Each ray shot through a pixel during one iteration of the path tracing algorithm is jittered, or moved ever-so-slightly, within the bounds of that pixel in relation to the ray sampled for that pixel during the previous iteration, and then the results for every iteration are averaged together. This results in edge pixels getting smoothed with their neighboring pixels which softens hard edges. + +I say supersampled anti-aliasing comes for free to a path tracer, because very little overhead is required to take advantage of it. This is in contrast to other renderers, such as ray tracers, where only a single ray is required for each pixel to generate a complete image. For a ray tracer to take advantage of supersampled anti-aliasing, additional rays need to be generated for each pixel that are not needed for basic image generation. As a result, supersampled anti-aliasing in a ray tracer increases runtime considerably. + +In my implementation, I discretize each pixel into a pre-defined number of rows and columns, and at each iteration I sample a pixel at a random location within one of the cells formed by the grid of rows and columns. The sub-pixel cell I sample changes each iteration, and I cycle through the cells linearly. The results of my implementation can be seen in the two images below. The first image does not use anti-aliasing. The second image does. The benefits of anti-aliasing in these images are most evident in the base of the green wall and in the outline of the red sphere against the white wall behind it. + +![alt tag](https://raw.githubusercontent.com/drerucha/Project3-Pathtracer/master/data/readme_pics/aa_without.jpg) + +![alt tag](https://raw.githubusercontent.com/drerucha/Project3-Pathtracer/master/data/readme_pics/aa_with.jpg) + +## PERFORMANCE ANALYSIS + +Below, I ran a few performance tests to measure the usefulness of stream compaction in my implementation. The tests were performed over 100 iterations with a block size locked at 128 threads-per-block with a varying trace depth (the number of times a ray was allowed to bounce in the scene). As can be seen in the scatter plot, no stream compaction outperforms stream compaction until a trace depth of 7-8 is reached. After that, stream compaction far outperforms no stream compaction. + +Considering what I wrote above about how stream compaction in concert with a per-ray parallelization scheme prevents rays of wildly varying "number-of-trace-depths-until-retirement" to share a warp on the GPU (see section labeled "Parallelization scheme" above for more information), these results make sense. With a low trace depth, the largest disparity between rays that retire at a trace depth of one vs. those that retire at the maximum trace depth is relatively small. As the maximum trace depth grows, this potential disparity grows as well, which results in more wasted GPU resources as kernel calls become filled with retired rays that can no longer contribute to the final render. + +The second chart visualizes the optimum block size for my path tracer. As with the previous tests, these tests were performed over 100 iterations. As can be seen in the chart, a block size of 64 or 128 is recommended for highest performance. + +![alt tag](https://raw.githubusercontent.com/drerucha/Project3-Pathtracer/master/data/readme_pics/chart_stream_compaction_performance.jpg) + +![alt tag](https://raw.githubusercontent.com/drerucha/Project3-Pathtracer/master/data/readme_pics/chart_block_size_comparison.jpg) + +## FUTURE WORK + +My next steps for this project include implementing: +* Direct lighting samples so my renders converge more efficiently. +* Bump maps. +* Image-based emittance. +* Depth of field. +* An obj loader. + +## FUN + +During development, I noticed that restricting random number generation in nonsensical ways resulted in some interesting abstract image creation. The first image below is no way physically correct, but it is my favorite image generated with my path tracer so far. + +![alt tag](https://raw.githubusercontent.com/drerucha/Project3-Pathtracer/master/data/readme_pics/fun_01.jpg) + +![alt tag](https://raw.githubusercontent.com/drerucha/Project3-Pathtracer/master/data/readme_pics/fun_02.jpg) + +## SPECIAL THANKS + +I want to give a quick shout-out to Patrick Cozzi who led the fall 2014 CIS 565 course at Penn, Harmony Li who was the TA for the same course, and Yining Karl Li who constructed much of the framework my path tracer was built upon. Thanks guys! \ No newline at end of file diff --git a/data/readme_pics/aa_with.jpg b/data/readme_pics/aa_with.jpg new file mode 100644 index 0000000..ac66439 Binary files /dev/null and b/data/readme_pics/aa_with.jpg differ diff --git a/data/readme_pics/aa_without.jpg b/data/readme_pics/aa_without.jpg new file mode 100644 index 0000000..bb8636e Binary files /dev/null and b/data/readme_pics/aa_without.jpg differ diff --git a/data/readme_pics/chart_block_size_comparison.jpg b/data/readme_pics/chart_block_size_comparison.jpg new file mode 100644 index 0000000..90fcf33 Binary files /dev/null and b/data/readme_pics/chart_block_size_comparison.jpg differ diff --git a/data/readme_pics/chart_stream_compaction_performance.jpg b/data/readme_pics/chart_stream_compaction_performance.jpg new file mode 100644 index 0000000..c7c7720 Binary files /dev/null and b/data/readme_pics/chart_stream_compaction_performance.jpg differ diff --git a/data/readme_pics/fresnel_refractions.jpg b/data/readme_pics/fresnel_refractions.jpg new file mode 100644 index 0000000..f0b8fe4 Binary files /dev/null and b/data/readme_pics/fresnel_refractions.jpg differ diff --git a/data/readme_pics/fun_01.jpg b/data/readme_pics/fun_01.jpg new file mode 100644 index 0000000..93b5330 Binary files /dev/null and b/data/readme_pics/fun_01.jpg differ diff --git a/data/readme_pics/fun_02.jpg b/data/readme_pics/fun_02.jpg new file mode 100644 index 0000000..f48f6f6 Binary files /dev/null and b/data/readme_pics/fun_02.jpg differ diff --git a/data/readme_pics/texture_mapping.jpg b/data/readme_pics/texture_mapping.jpg new file mode 100644 index 0000000..c62280c Binary files /dev/null and b/data/readme_pics/texture_mapping.jpg differ diff --git a/data/readme_pics/texture_mapping_02.jpg b/data/readme_pics/texture_mapping_02.jpg new file mode 100644 index 0000000..51c760c Binary files /dev/null and b/data/readme_pics/texture_mapping_02.jpg differ diff --git a/data/scenes/sampleScene_02.txt b/data/scenes/sampleScene_02.txt new file mode 100644 index 0000000..b595bd6 --- /dev/null +++ b/data/scenes/sampleScene_02.txt @@ -0,0 +1,225 @@ +MATERIAL 0 //white diffuse +RGB 1 1 1 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 0 +REFRIOR 0 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 0 + +MATERIAL 1 //red diffuse +RGB .63 .06 .04 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 0 +REFRIOR 0 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 0 + +MATERIAL 2 //green diffuse +RGB .15 .48 .09 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 0 +REFRIOR 0 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 0 + +MATERIAL 3 //red glossy +RGB .63 .06 .04 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 0 +REFRIOR 2 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 0 + +MATERIAL 4 //white glossy +RGB 1 1 1 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 0 +REFRIOR 2 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 0 + +MATERIAL 5 //glass +RGB 0 0 0 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 1 +REFRIOR 2.2 +SCATTER 0 +ABSCOEFF .02 5.1 5.7 +RSCTCOEFF 13 +EMITTANCE 0 + +MATERIAL 6 //green glossy +RGB .15 .48 .09 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 0 +REFRIOR 2.6 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 0 + +MATERIAL 7 //light +RGB 1 1 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 1 + +MATERIAL 8 //light +RGB 1 1 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 15 + +TEXTURE 0 +FILE world_map_256x128.bmp + +TEXTURE 1 +FILE stripes_256x128.bmp + +TEXTURE 2 +FILE black_and_white_tiles_256x256.bmp + +CAMERA +RES 800 800 +FOVY 25 +ITERATIONS 5000 +FILE test.bmp +frame 0 +EYE 0 4.5 12 +VIEW 0 0 -1 +UP 0 1 0 + +OBJECT 0 +cube +material 0 +texture 2 +frame 0 +TRANS 0 0 0 +ROTAT 0 0 90 +SCALE .01 10 10 + +OBJECT 1 +cube +material 0 +texture 1 +frame 0 +TRANS 0 5 -5 +ROTAT 0 90 0 +SCALE .01 10 10 + +OBJECT 2 +cube +material 0 +texture -1 +frame 0 +TRANS 0 10 0 +ROTAT 0 0 90 +SCALE .01 10 10 + +OBJECT 3 +cube +material 1 +texture -1 +frame 0 +TRANS -5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +OBJECT 4 +cube +material 2 +texture -1 +frame 0 +TRANS 5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +OBJECT 5 +sphere +material 5 +texture -1 +frame 0 +TRANS 0 2 0 +ROTAT 0 90 0 +SCALE 3 3 3 + +OBJECT 6 +sphere +material 3 +texture -1 +frame 0 +TRANS 2 5 2 +ROTAT 0 0 0 +SCALE 2.5 2.5 2.5 + +OBJECT 7 +sphere +material 3 +texture 0 +frame 0 +TRANS -2 5 -2 +ROTAT 0 180 0 +SCALE 3 3 3 + +OBJECT 8 +cube +material 8 +texture -1 +frame 0 +TRANS 0 10 0 +ROTAT 0 0 90 +SCALE .3 3 3 + +//OBJECT 9 +//sphere +//material 8 +//texture -1 +//frame 0 +//TRANS 2 2 16 +//ROTAT 0 180 0 +//SCALE 1 1 1 + +//OBJECT 10 +//cube +//material 0 +//texture -1 +//frame 0 +//TRANS 0 5 5 +//ROTAT 0 90 0 +//SCALE .01 10 10 \ No newline at end of file diff --git a/data/textures/black_and_white_tiles.jpg b/data/textures/black_and_white_tiles.jpg new file mode 100644 index 0000000..aa90c4d Binary files /dev/null and b/data/textures/black_and_white_tiles.jpg differ diff --git a/data/textures/black_and_white_tiles_256x256.bmp b/data/textures/black_and_white_tiles_256x256.bmp new file mode 100644 index 0000000..03a3693 Binary files /dev/null and b/data/textures/black_and_white_tiles_256x256.bmp differ diff --git a/data/textures/stripes.jpg b/data/textures/stripes.jpg new file mode 100644 index 0000000..1dcd7ef Binary files /dev/null and b/data/textures/stripes.jpg differ diff --git a/data/textures/stripes_256x128.bmp b/data/textures/stripes_256x128.bmp new file mode 100644 index 0000000..007f88f Binary files /dev/null and b/data/textures/stripes_256x128.bmp differ diff --git a/data/textures/world_map_1024x512.bmp b/data/textures/world_map_1024x512.bmp new file mode 100644 index 0000000..058dcec Binary files /dev/null and b/data/textures/world_map_1024x512.bmp differ diff --git a/data/textures/world_map_128x64.bmp b/data/textures/world_map_128x64.bmp new file mode 100644 index 0000000..6630a8a Binary files /dev/null and b/data/textures/world_map_128x64.bmp differ diff --git a/data/textures/world_map_256x128.bmp b/data/textures/world_map_256x128.bmp new file mode 100644 index 0000000..416f3e9 Binary files /dev/null and b/data/textures/world_map_256x128.bmp differ diff --git a/renders/bmp/2014-10-02_01_starting_point.bmp b/renders/bmp/2014-10-02_01_starting_point.bmp new file mode 100644 index 0000000..b701a7c Binary files /dev/null and b/renders/bmp/2014-10-02_01_starting_point.bmp differ diff --git a/renders/bmp/2014-10-03_01_ray_direction_test.bmp b/renders/bmp/2014-10-03_01_ray_direction_test.bmp new file mode 100644 index 0000000..d9004cf Binary files /dev/null and b/renders/bmp/2014-10-03_01_ray_direction_test.bmp differ diff --git a/renders/bmp/2014-10-03_02_sphere_intersection_test_01.bmp b/renders/bmp/2014-10-03_02_sphere_intersection_test_01.bmp new file mode 100644 index 0000000..cac22b4 Binary files /dev/null and b/renders/bmp/2014-10-03_02_sphere_intersection_test_01.bmp differ diff --git a/renders/bmp/2014-10-03_03_sphere_intersection_test_02.bmp b/renders/bmp/2014-10-03_03_sphere_intersection_test_02.bmp new file mode 100644 index 0000000..cfb1f02 Binary files /dev/null and b/renders/bmp/2014-10-03_03_sphere_intersection_test_02.bmp differ diff --git a/renders/bmp/2014-10-03_04_add_material_references.bmp b/renders/bmp/2014-10-03_04_add_material_references.bmp new file mode 100644 index 0000000..6ec42be Binary files /dev/null and b/renders/bmp/2014-10-03_04_add_material_references.bmp differ diff --git a/renders/bmp/2014-10-03_05_box_intersections.bmp b/renders/bmp/2014-10-03_05_box_intersections.bmp new file mode 100644 index 0000000..778c258 Binary files /dev/null and b/renders/bmp/2014-10-03_05_box_intersections.bmp differ diff --git a/renders/bmp/2014-10-03_06_incorrect_intersection_normals.bmp b/renders/bmp/2014-10-03_06_incorrect_intersection_normals.bmp new file mode 100644 index 0000000..e7e57fb Binary files /dev/null and b/renders/bmp/2014-10-03_06_incorrect_intersection_normals.bmp differ diff --git a/renders/bmp/2014-10-03_07_fixed_intersection_normals.bmp b/renders/bmp/2014-10-03_07_fixed_intersection_normals.bmp new file mode 100644 index 0000000..4ba5eae Binary files /dev/null and b/renders/bmp/2014-10-03_07_fixed_intersection_normals.bmp differ diff --git a/renders/bmp/2014-10-05_01_first_image.bmp b/renders/bmp/2014-10-05_01_first_image.bmp new file mode 100644 index 0000000..ab136cf Binary files /dev/null and b/renders/bmp/2014-10-05_01_first_image.bmp differ diff --git a/renders/bmp/2014-10-05_02_second_image.bmp b/renders/bmp/2014-10-05_02_second_image.bmp new file mode 100644 index 0000000..8ff1fe2 Binary files /dev/null and b/renders/bmp/2014-10-05_02_second_image.bmp differ diff --git a/renders/bmp/2014-10-05_03_perfect_specular.bmp b/renders/bmp/2014-10-05_03_perfect_specular.bmp new file mode 100644 index 0000000..f68700b Binary files /dev/null and b/renders/bmp/2014-10-05_03_perfect_specular.bmp differ diff --git a/renders/bmp/2014-10-09_01_first_refraction_attempt.bmp b/renders/bmp/2014-10-09_01_first_refraction_attempt.bmp new file mode 100644 index 0000000..657c167 Binary files /dev/null and b/renders/bmp/2014-10-09_01_first_refraction_attempt.bmp differ diff --git a/renders/bmp/2014-10-09_02_second_refraction_attempt.bmp b/renders/bmp/2014-10-09_02_second_refraction_attempt.bmp new file mode 100644 index 0000000..4218c03 Binary files /dev/null and b/renders/bmp/2014-10-09_02_second_refraction_attempt.bmp differ diff --git a/renders/bmp/2014-10-09_03_diffuse_with_aa.bmp b/renders/bmp/2014-10-09_03_diffuse_with_aa.bmp new file mode 100644 index 0000000..7674334 Binary files /dev/null and b/renders/bmp/2014-10-09_03_diffuse_with_aa.bmp differ diff --git a/renders/bmp/2014-10-09_04_glass_with_aa.bmp b/renders/bmp/2014-10-09_04_glass_with_aa.bmp new file mode 100644 index 0000000..5ccc9a0 Binary files /dev/null and b/renders/bmp/2014-10-09_04_glass_with_aa.bmp differ diff --git a/renders/bmp/2014-10-11_01_texture_mapping_01.bmp b/renders/bmp/2014-10-11_01_texture_mapping_01.bmp new file mode 100644 index 0000000..c1fa08a Binary files /dev/null and b/renders/bmp/2014-10-11_01_texture_mapping_01.bmp differ diff --git a/renders/bmp/2014-10-11_02_texture_mapping_02.bmp b/renders/bmp/2014-10-11_02_texture_mapping_02.bmp new file mode 100644 index 0000000..9e53c4e Binary files /dev/null and b/renders/bmp/2014-10-11_02_texture_mapping_02.bmp differ diff --git a/renders/bmp/2014-10-11_03_texture_mapping_03.bmp b/renders/bmp/2014-10-11_03_texture_mapping_03.bmp new file mode 100644 index 0000000..1f5108c Binary files /dev/null and b/renders/bmp/2014-10-11_03_texture_mapping_03.bmp differ diff --git a/renders/bmp/2014-10-12_01_texture_mapping.bmp b/renders/bmp/2014-10-12_01_texture_mapping.bmp new file mode 100644 index 0000000..ad62919 Binary files /dev/null and b/renders/bmp/2014-10-12_01_texture_mapping.bmp differ diff --git a/renders/bmp/2014-10-12_02_cube_map_test.bmp b/renders/bmp/2014-10-12_02_cube_map_test.bmp new file mode 100644 index 0000000..a8988e7 Binary files /dev/null and b/renders/bmp/2014-10-12_02_cube_map_test.bmp differ diff --git a/renders/bmp/5000_iterations_bounded_random_number_generation.bmp b/renders/bmp/5000_iterations_bounded_random_number_generation.bmp new file mode 100644 index 0000000..0061235 Binary files /dev/null and b/renders/bmp/5000_iterations_bounded_random_number_generation.bmp differ diff --git a/renders/bmp/5000_iterations_bounded_random_number_generation_02.bmp b/renders/bmp/5000_iterations_bounded_random_number_generation_02.bmp new file mode 100644 index 0000000..2e4cbb4 Binary files /dev/null and b/renders/bmp/5000_iterations_bounded_random_number_generation_02.bmp differ diff --git a/renders/bmp/5000_iterations_unbounded_random_number_generation.bmp b/renders/bmp/5000_iterations_unbounded_random_number_generation.bmp new file mode 100644 index 0000000..c7caddc Binary files /dev/null and b/renders/bmp/5000_iterations_unbounded_random_number_generation.bmp differ diff --git a/renders/bmp/random_numbers_mod_10.bmp b/renders/bmp/random_numbers_mod_10.bmp new file mode 100644 index 0000000..e3a783a Binary files /dev/null and b/renders/bmp/random_numbers_mod_10.bmp differ diff --git a/renders/bmp/random_numbers_mod_100.bmp b/renders/bmp/random_numbers_mod_100.bmp new file mode 100644 index 0000000..6c35e6a Binary files /dev/null and b/renders/bmp/random_numbers_mod_100.bmp differ diff --git a/renders/bmp/random_numbers_mod_1000.bmp b/renders/bmp/random_numbers_mod_1000.bmp new file mode 100644 index 0000000..f3a2175 Binary files /dev/null and b/renders/bmp/random_numbers_mod_1000.bmp differ diff --git a/renders/bmp/random_numbers_mod_10000.bmp b/renders/bmp/random_numbers_mod_10000.bmp new file mode 100644 index 0000000..3880500 Binary files /dev/null and b/renders/bmp/random_numbers_mod_10000.bmp differ diff --git a/renders/bmp/random_numbers_mod_100000.bmp b/renders/bmp/random_numbers_mod_100000.bmp new file mode 100644 index 0000000..a8544a2 Binary files /dev/null and b/renders/bmp/random_numbers_mod_100000.bmp differ diff --git a/renders/bmp/random_numbers_mod_1000000.bmp b/renders/bmp/random_numbers_mod_1000000.bmp new file mode 100644 index 0000000..30b1ec0 Binary files /dev/null and b/renders/bmp/random_numbers_mod_1000000.bmp differ diff --git a/renders/jpg/2014-10-02_01_starting_point.jpg b/renders/jpg/2014-10-02_01_starting_point.jpg new file mode 100644 index 0000000..068c433 Binary files /dev/null and b/renders/jpg/2014-10-02_01_starting_point.jpg differ diff --git a/renders/jpg/2014-10-03_01_ray_direction_test.jpg b/renders/jpg/2014-10-03_01_ray_direction_test.jpg new file mode 100644 index 0000000..b420a67 Binary files /dev/null and b/renders/jpg/2014-10-03_01_ray_direction_test.jpg differ diff --git a/renders/jpg/2014-10-03_02_sphere_intersection_test_01.jpg b/renders/jpg/2014-10-03_02_sphere_intersection_test_01.jpg new file mode 100644 index 0000000..65e77ef Binary files /dev/null and b/renders/jpg/2014-10-03_02_sphere_intersection_test_01.jpg differ diff --git a/renders/jpg/2014-10-03_03_sphere_intersection_test_02.jpg b/renders/jpg/2014-10-03_03_sphere_intersection_test_02.jpg new file mode 100644 index 0000000..ad3e285 Binary files /dev/null and b/renders/jpg/2014-10-03_03_sphere_intersection_test_02.jpg differ diff --git a/renders/jpg/2014-10-03_04_add_material_references.jpg b/renders/jpg/2014-10-03_04_add_material_references.jpg new file mode 100644 index 0000000..b217fb6 Binary files /dev/null and b/renders/jpg/2014-10-03_04_add_material_references.jpg differ diff --git a/renders/jpg/2014-10-03_05_box_intersections.jpg b/renders/jpg/2014-10-03_05_box_intersections.jpg new file mode 100644 index 0000000..cf65f2c Binary files /dev/null and b/renders/jpg/2014-10-03_05_box_intersections.jpg differ diff --git a/renders/jpg/2014-10-03_06_incorrect_intersection_normals.jpg b/renders/jpg/2014-10-03_06_incorrect_intersection_normals.jpg new file mode 100644 index 0000000..a41d60b Binary files /dev/null and b/renders/jpg/2014-10-03_06_incorrect_intersection_normals.jpg differ diff --git a/renders/jpg/2014-10-03_07_fixed_intersection_normals.jpg b/renders/jpg/2014-10-03_07_fixed_intersection_normals.jpg new file mode 100644 index 0000000..b0aaedd Binary files /dev/null and b/renders/jpg/2014-10-03_07_fixed_intersection_normals.jpg differ diff --git a/renders/jpg/2014-10-05_01_first_image.jpg b/renders/jpg/2014-10-05_01_first_image.jpg new file mode 100644 index 0000000..dfa60ee Binary files /dev/null and b/renders/jpg/2014-10-05_01_first_image.jpg differ diff --git a/renders/jpg/2014-10-05_02_second_image.jpg b/renders/jpg/2014-10-05_02_second_image.jpg new file mode 100644 index 0000000..bb8636e Binary files /dev/null and b/renders/jpg/2014-10-05_02_second_image.jpg differ diff --git a/renders/jpg/2014-10-05_03_perfect_specular.jpg b/renders/jpg/2014-10-05_03_perfect_specular.jpg new file mode 100644 index 0000000..d8636d0 Binary files /dev/null and b/renders/jpg/2014-10-05_03_perfect_specular.jpg differ diff --git a/renders/jpg/2014-10-09_01_first_refraction_attempt.jpg b/renders/jpg/2014-10-09_01_first_refraction_attempt.jpg new file mode 100644 index 0000000..e8fa777 Binary files /dev/null and b/renders/jpg/2014-10-09_01_first_refraction_attempt.jpg differ diff --git a/renders/jpg/2014-10-11_01_texture_mapping_01.jpg b/renders/jpg/2014-10-11_01_texture_mapping_01.jpg new file mode 100644 index 0000000..3d3650c Binary files /dev/null and b/renders/jpg/2014-10-11_01_texture_mapping_01.jpg differ diff --git a/renders/jpg/2014-10-12_01_texture_mapping.jpg b/renders/jpg/2014-10-12_01_texture_mapping.jpg new file mode 100644 index 0000000..a20618b Binary files /dev/null and b/renders/jpg/2014-10-12_01_texture_mapping.jpg differ diff --git a/renders/jpg/2014-10-12_02_cube_map_test.jpg b/renders/jpg/2014-10-12_02_cube_map_test.jpg new file mode 100644 index 0000000..c62280c Binary files /dev/null and b/renders/jpg/2014-10-12_02_cube_map_test.jpg differ diff --git a/src/SimpleTimer.cpp b/src/SimpleTimer.cpp new file mode 100644 index 0000000..03b07ae --- /dev/null +++ b/src/SimpleTimer.cpp @@ -0,0 +1,44 @@ +#include "SimpleTimer.h" + + +SimpleTimer::SimpleTimer() +{ +} + + +SimpleTimer::~SimpleTimer() +{ +} + + +void SimpleTimer::start() +{ + start_time = GetTimeMs64(); +} + + +float SimpleTimer::stop() +{ + __int64 end_time = GetTimeMs64(); + return ( float )( end_time - start_time ); +} + + +__int64 SimpleTimer::GetTimeMs64() +{ + /* Windows */ + FILETIME ft; + LARGE_INTEGER li; + + /* Get the amount of 100 nano seconds intervals elapsed since January 1, 1601 (UTC) and copy it + * to a LARGE_INTEGER structure. */ + GetSystemTimeAsFileTime(&ft); + li.LowPart = ft.dwLowDateTime; + li.HighPart = ft.dwHighDateTime; + + __int64 ret = li.QuadPart; + ret -= 116444736000000000LL; /* Convert from file time to UNIX epoch time. */ + ret /= 10000; /* From 100 nano seconds (10^-7) to 1 millisecond (10^-3) intervals */ + + return ret; +} \ No newline at end of file diff --git a/src/SimpleTimer.h b/src/SimpleTimer.h new file mode 100644 index 0000000..a2a51b0 --- /dev/null +++ b/src/SimpleTimer.h @@ -0,0 +1,23 @@ +#pragma once + +#ifndef _SIMPLE_TIMER +#define _SIMPLE_TIMER + +#include + +class SimpleTimer +{ +public: + SimpleTimer( void ); + ~SimpleTimer( void ); + + void start( void ); + float stop( void ); + +private: + __int64 GetTimeMs64( void ); + + __int64 start_time; +}; + +#endif \ No newline at end of file diff --git a/src/image.cpp b/src/image.cpp index 1e0e072..655808c 100644 --- a/src/image.cpp +++ b/src/image.cpp @@ -12,10 +12,12 @@ image::image(int x, int y){ xSize = x; ySize = y; + redChannel = new float[x*y]; greenChannel = new float[x*y]; blueChannel = new float[x*y]; alphaChannel = new float[x*y]; + for(int i=0; i<(x*y); i++){ redChannel[i] = 0; greenChannel[i] = 0; diff --git a/src/image.h b/src/image.h index 29b8a62..5f51608 100644 --- a/src/image.h +++ b/src/image.h @@ -7,6 +7,7 @@ #define Raytracer_image_h #include "glm/glm.hpp" +//#include using namespace std; @@ -26,13 +27,12 @@ class image{ int ySize; gammaSettings gamma; public: - image(int x, int y); + image(int x=0, int y=0); ~image(); //------------------------ //-------GETTERS---------- //------------------------ - glm::vec3 readPixelRGB(int x, int y); glm::vec4 readPixelRGBA(int x, int y); float readPixelR(int x, int y); float readPixelG(int x, int y); @@ -44,8 +44,11 @@ class image{ float* getAlphaChannel(); glm::vec3* getRGBChannels(); glm::vec4* getRGBAChannels(); - glm::vec2 getDimensions(); gammaSettings getGammaSettings(); + + glm::vec3 readPixelRGB(int x, int y); + glm::vec2 getDimensions(); + //------------------------ //-------SETTERS---------- diff --git a/src/interactions.h b/src/interactions.h index 7bf6fab..da40ec1 100644 --- a/src/interactions.h +++ b/src/interactions.h @@ -18,6 +18,7 @@ struct AbsorptionAndScatteringProperties{ float reducedScatteringCoefficient; }; + // Forward declaration __host__ __device__ bool calculateScatterAndAbsorption(ray& r, float& depth, AbsorptionAndScatteringProperties& currentAbsorptionAndScattering, glm::vec3& unabsorbedColor, material m, float randomFloatForScatteringDistance, float randomFloat2, float randomFloat3); __host__ __device__ glm::vec3 getRandomDirectionInSphere(float xi1, float xi2); @@ -27,80 +28,176 @@ __host__ __device__ glm::vec3 calculateReflectionDirection(glm::vec3 normal, glm __host__ __device__ Fresnel calculateFresnel(glm::vec3 normal, glm::vec3 incident, float incidentIOR, float transmittedIOR, glm::vec3 reflectionDirection, glm::vec3 transmissionDirection); __host__ __device__ glm::vec3 calculateRandomDirectionInHemisphere(glm::vec3 normal, float xi1, float xi2); + // TODO (OPTIONAL): IMPLEMENT THIS FUNCTION -__host__ __device__ glm::vec3 calculateTransmission(glm::vec3 absorptionCoefficient, float distance) { - return glm::vec3(0,0,0); +__host__ +__device__ +glm::vec3 calculateTransmission( glm::vec3 absorptionCoefficient, + float distance ) +{ + return glm::vec3( 0.0f, 0.0f, 0.0f ); } + // TODO (OPTIONAL): IMPLEMENT THIS FUNCTION -__host__ __device__ bool calculateScatterAndAbsorption(ray& r, float& depth, AbsorptionAndScatteringProperties& currentAbsorptionAndScattering, - glm::vec3& unabsorbedColor, material m, float randomFloatForScatteringDistance, float randomFloat2, float randomFloat3){ - return false; +__host__ +__device__ +bool calculateScatterAndAbsorption( ray& r, + float& depth, + AbsorptionAndScatteringProperties& currentAbsorptionAndScattering, + glm::vec3& unabsorbedColor, + material m, + float randomFloatForScatteringDistance, + float randomFloat2, + float randomFloat3 ) +{ + return false; } + // TODO (OPTIONAL): IMPLEMENT THIS FUNCTION -__host__ __device__ glm::vec3 calculateTransmissionDirection(glm::vec3 normal, glm::vec3 incident, float incidentIOR, float transmittedIOR) { - return glm::vec3(0,0,0); +__host__ +__device__ +glm::vec3 calculateTransmissionDirection( glm::vec3 normal, // Intersection normal. + glm::vec3 incident, // Incoming ray direction. + float ior_incident, // Incoming index of refraction. + float ior_transmitted ) // Outgoing index of refraction. +{ + glm::vec3 normal_oriented = ( glm::dot( normal, incident ) < 0.0f ) ? normal : ( -1.0f * normal ); + bool ray_is_entering = ( glm::dot( normal, normal_oriented ) > 0.0f ); + + float n = ior_incident / ior_transmitted; + float c1 = glm::dot( incident, normal_oriented ); + float c2 = 1.0f - ( n * n * ( 1.0f - c1 * c1 ) ); + + // If angle is too shallow, then all light is reflected. + if ( c2 < 0.0f ) { + //return glm::vec3( 0.0f, 0.0f, 0.0f ); + return calculateReflectionDirection( normal, incident ); + } + + return glm::normalize( ( incident * n ) - ( normal * ( ( ray_is_entering ? 1.0f : -1.0f ) * ( c1 * n + sqrt( c2 ) ) ) ) ); } + // TODO (OPTIONAL): IMPLEMENT THIS FUNCTION -__host__ __device__ glm::vec3 calculateReflectionDirection(glm::vec3 normal, glm::vec3 incident) { - //nothing fancy here - return glm::vec3(0,0,0); +__host__ +__device__ +glm::vec3 calculateReflectionDirection( glm::vec3 normal, // Intersection normal. + glm::vec3 incident ) // Incoming ray direction. +{ + //glm::vec3 normal_oriented = ( glm::dot( normal, incident ) < 0.0f ) ? normal : ( -1.0f * normal ); + //return glm::normalize( incident - ( 2.0f * glm::dot( normal_oriented, incident ) * normal_oriented ) ); + + return incident - ( 2.0f * glm::dot( normal, incident ) * normal ); } -// TODO (OPTIONAL): IMPLEMENT THIS FUNCTION -__host__ __device__ Fresnel calculateFresnel(glm::vec3 normal, glm::vec3 incident, float incidentIOR, float transmittedIOR, glm::vec3 reflectionDirection, glm::vec3 transmissionDirection) { - Fresnel fresnel; - fresnel.reflectionCoefficient = 1; - fresnel.transmissionCoefficient = 0; - return fresnel; +// TODO (OPTIONAL): IMPLEMENT THIS FUNCTION +__host__ +__device__ +Fresnel calculateFresnel( glm::vec3 normal, // Intersection normal. + glm::vec3 incident, // Incoming ray direction. + float ior_incident, // Incoming index of refraction. + float ior_transmitted, // Outgoing index of refraction. + glm::vec3 reflection_dir, // Reflecton direction. + glm::vec3 transmission_dir ) // Transmission direction. +{ + float n = ior_incident / ior_transmitted; + float a = n - 1.0f; + float b = n + 1.0f; + float Ro = ( a * a ) / ( b * b ); + + glm::vec3 normal_oriented = ( glm::dot( normal, incident ) < 0.0f ) ? normal : ( -1.0f * normal ); + bool ray_is_entering = ( glm::dot( normal, normal_oriented ) > 0.0f ); + float angle = ray_is_entering ? glm::dot( -incident, normal ) : glm::dot( transmission_dir, normal ); + float c = 1.0f - angle; + + float R = Ro + ( 1.0f - Ro ) * c * c * c * c * c; + + //float P = 0.25f + 0.5f * R; + //float RP = R / P; + //float TP = ( 1.0f - R ) / ( 1.0f - P ); + + Fresnel fresnel; + fresnel.reflectionCoefficient = R; + fresnel.transmissionCoefficient = 1.0f - R; + + return fresnel; } + // LOOK: This function demonstrates cosine weighted random direction generation in a sphere! -__host__ __device__ glm::vec3 calculateRandomDirectionInHemisphere(glm::vec3 normal, float xi1, float xi2) { - - // Crucial difference between this and calculateRandomDirectionInSphere: THIS IS COSINE WEIGHTED! - - float up = sqrt(xi1); // cos(theta) - float over = sqrt(1 - up * up); // sin(theta) - float around = xi2 * TWO_PI; +__host__ +__device__ +glm::vec3 calculateRandomDirectionInHemisphere( glm::vec3 normal, + float xi1, + float xi2 ) +{ + // Crucial difference between this and calculateRandomDirectionInSphere: THIS IS COSINE WEIGHTED! - // 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. + float up = sqrt( xi1 ); // cos(theta) + float over = sqrt( 1.0f - up * up ); // sin(theta) + float around = xi2 * TWO_PI; - 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); - } + // 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. - // 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)); + glm::vec3 directionNotNormal; + if ( abs( normal.x ) < SQRT_OF_ONE_THIRD ) { + directionNotNormal = glm::vec3( 1.0f, 0.0f, 0.0f ); + } else if ( abs( normal.y ) < SQRT_OF_ONE_THIRD ) { + directionNotNormal = glm::vec3( 0.0f, 1.0f, 0.0f ); + } else { + directionNotNormal = glm::vec3( 0.0f, 0.0f, 1.0f ); + } - return ( up * normal ) + ( cos(around) * over * perpendicularDirection1 ) + ( sin(around) * over * perpendicularDirection2 ); + // 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 ); } + // TODO: IMPLEMENT THIS FUNCTION // Now that you know how cosine weighted direction generation works, try implementing // non-cosine (uniform) weighted random direction generation. // This should be much easier than if you had to implement calculateRandomDirectionInHemisphere. -__host__ __device__ glm::vec3 getRandomDirectionInSphere(float xi1, float xi2) { - return glm::vec3(0,0,0); +__host__ +__device__ +glm::vec3 getRandomDirectionInSphere( float xi1, + float xi2 ) +{ + float r = sqrt( 1.0f - xi1 * xi1 ); // sin(theta) + float phi = xi2 * TWO_PI; + return glm::vec3( r * cos( phi ), r * sin( phi ), xi1 ); } + // TODO (PARTIALLY OPTIONAL): IMPLEMENT THIS FUNCTION // Returns 0 if diffuse scatter, 1 if reflected, 2 if transmitted. -__host__ __device__ int calculateBSDF(ray& r, glm::vec3 intersect, glm::vec3 normal, glm::vec3 emittedColor, - AbsorptionAndScatteringProperties& currentAbsorptionAndScattering, - glm::vec3& color, glm::vec3& unabsorbedColor, material m){ - - return 1; +__host__ +__device__ +int calculateBSDF( ray& r, + glm::vec3 intersect, + glm::vec3 normal, + glm::vec3 emittedColor, + AbsorptionAndScatteringProperties& currentAbsorptionAndScattering, + glm::vec3& color, + glm::vec3& unabsorbedColor, + material m ) +{ + //// Diffuse. + //if ( !m.hasReflective && !m.hasRefractive ) { + // return 0; + //} + //// Perfect specular. + //else { + // return 1; + //} + + return 0; }; #endif diff --git a/src/intersections.h b/src/intersections.h index c9eafb6..e9cd601 100644 --- a/src/intersections.h +++ b/src/intersections.h @@ -16,14 +16,18 @@ // Some forward declarations __host__ __device__ glm::vec3 getPointOnRay(ray r, float t); __host__ __device__ glm::vec3 multiplyMV(cudaMat4 m, glm::vec4 v); +__host__ __device__ cudaMat4 transposeM( cudaMat4 m ); __host__ __device__ glm::vec3 getSignOfRay(ray r); __host__ __device__ glm::vec3 getInverseDirectionOfRay(ray r); +__host__ __device__ void swapFloats( float &a, float &b ); __host__ __device__ float boxIntersectionTest(staticGeom sphere, ray r, glm::vec3& intersectionPoint, glm::vec3& normal); __host__ __device__ float sphereIntersectionTest(staticGeom sphere, ray r, glm::vec3& intersectionPoint, glm::vec3& normal); __host__ __device__ glm::vec3 getRandomPointOnCube(staticGeom cube, float randomSeed); +__host__ __device__ glm::vec3 computeBoxObjectSpaceNormal( const glm::vec3 &object_space_intersection_point ); + // Handy dandy little hashing function that provides seeds for random number generation -__host__ __device__ unsigned int hash(unsigned int a){ +__host__ __device__ unsigned int simpleHash(unsigned int a){ a = (a+0x7ed55d16) + (a<<12); a = (a^0xc761c23c) ^ (a>>19); a = (a+0x165667b1) + (a<<5); @@ -33,6 +37,7 @@ __host__ __device__ unsigned int hash(unsigned int a){ return a; } + // Quick and dirty epsilon check __host__ __device__ bool epsilonCheck(float a, float b){ if(fabs(fabs(a)-fabs(b)) < EPSILON){ @@ -42,11 +47,13 @@ __host__ __device__ bool epsilonCheck(float a, float b){ } } + // Self explanatory __host__ __device__ glm::vec3 getPointOnRay(ray r, float t){ return r.origin + float(t - .0001f) * glm::normalize(r.direction); } + // LOOK: This is a custom function for multiplying cudaMat4 4x4 matrixes with vectors. // This is a workaround for GLM matrix multiplication not working properly on pre-Fermi NVIDIA GPUs. // Multiplies a cudaMat4 matrix and a vec4 and returns a vec3 clipped from the vec4 @@ -58,64 +65,230 @@ __host__ __device__ glm::vec3 multiplyMV(cudaMat4 m, glm::vec4 v){ return r; } + +__host__ +__device__ +cudaMat4 transposeM( cudaMat4 m ) +{ + cudaMat4 transpose; + transpose.x = glm::vec4( m.x.x, m.y.x, m.z.x, m.w.x ); + transpose.y = glm::vec4( m.x.y, m.y.y, m.z.y, m.w.y ); + transpose.z = glm::vec4( m.x.z, m.y.z, m.z.z, m.w.z ); + transpose.w = glm::vec4( m.x.w, m.y.w, m.z.w, m.w.w ); + return transpose; +} + + // Gets 1/direction for a ray __host__ __device__ glm::vec3 getInverseDirectionOfRay(ray r){ return glm::vec3(1.0/r.direction.x, 1.0/r.direction.y, 1.0/r.direction.z); } + // Gets sign of each component of a ray's inverse direction __host__ __device__ glm::vec3 getSignOfRay(ray r){ glm::vec3 inv_direction = getInverseDirectionOfRay(r); return glm::vec3((int)(inv_direction.x < 0), (int)(inv_direction.y < 0), (int)(inv_direction.z < 0)); } + +// A swap method I can call from the GPU. +__host__ +__device__ +void swapFloats( float &a, float &b ) +{ + float c = a; + a = b; + b = c; +} + + // TODO: IMPLEMENT THIS FUNCTION // Cube intersection test, return -1 if no intersection, otherwise, distance to intersection -__host__ __device__ float boxIntersectionTest(staticGeom box, ray r, glm::vec3& intersectionPoint, glm::vec3& normal){ +__host__ +__device__ +float boxIntersectionTest( staticGeom box, + ray r, + glm::vec3& intersectionPoint, + glm::vec3& normal ) +{ + // Solution adapted from "Real-Time Rendering Third Edition" pg 742-744 by Akenine-Moller, Haines, and Hoffman. + // Slab method for AABB and OBB. + // Method assumes base side length of cube is 1. + + // Transform ray origin and direction to object space. + glm::vec3 ro = multiplyMV( box.inverseTransform, glm::vec4( r.origin, 1.0f ) ); + glm::vec3 rd = glm::normalize( multiplyMV( box.inverseTransform, glm::vec4( r.direction, 0.0f ) ) ); + + ray rt; + rt.origin = ro; + rt.direction = rd; + + // Cube center. + glm::vec3 ac( 0.0f, 0.0f, 0.0f ); + + // Normalized side vectors of cube. + glm::vec3 ax( 1.0f, 0.0f, 0.0f ); + glm::vec3 ay( 0.0f, 1.0f, 0.0f ); + glm::vec3 az( 0.0f, 0.0f, 1.0f ); - return -1; + // Cube half lengths. + float hx = 0.5f; + float hy = 0.5f; + float hz = 0.5f; + + glm::vec3 side_vectors[3] = { ax, ay, az }; + float half_lengths[3] = { hx, hy, hz }; + + float tmin = -FLT_MAX; + float tmax = FLT_MAX; + + glm::vec3 p = ac - rt.origin; + + for ( int i = 0; i < 3; ++i ) { + float e = glm::dot( side_vectors[i], p ); + float f = glm::dot( side_vectors[i], rt.direction ); + + // Ray is not parallel to current slab. + if ( abs( f ) > EPSILON ) { + float t1 = ( ( e + half_lengths[i] ) / f ); + float t2 = ( ( e - half_lengths[i] ) / f ); + + if ( t1 > t2 ) { + swapFloats( t1, t2 ); + } + + if ( t1 > tmin ) { + tmin = t1; + } + + if ( t2 < tmax ) { + tmax = t2; + } + + // Ray misses box, so reject. + if ( tmin > tmax ) { + return -1.0f; + } + + // Box is behind ray origin, so reject. + if ( tmax < 0.0f ) { + return -1.0f; + } + } + // Ray is parallel to current slab and ray is outside current slab, so reject. + else if ( -e - half_lengths[i] > 0.0f || -e + half_lengths[i] < 0.0f ) { + return -1.0f; + } + } + + glm::vec3 object_space_intersection_point; + glm::vec3 realIntersectionPoint; + if ( tmin > 0.0f ) { + object_space_intersection_point = getPointOnRay( rt, tmin ); + realIntersectionPoint = multiplyMV( box.transform, glm::vec4( object_space_intersection_point, 1.0f ) ); + } + else { + object_space_intersection_point = getPointOnRay( rt, tmax ); + realIntersectionPoint = multiplyMV( box.transform, glm::vec4( object_space_intersection_point, 1.0f ) ); + } + + intersectionPoint = realIntersectionPoint; + + glm::vec3 object_space_normal = computeBoxObjectSpaceNormal( object_space_intersection_point ); + normal = glm::normalize( multiplyMV( transposeM( box.inverseTransform ), glm::vec4( object_space_normal, 0.0f ) ) ); + + return glm::length( r.origin - realIntersectionPoint ); } + +__host__ +__device__ +glm::vec3 computeBoxObjectSpaceNormal( const glm::vec3 &object_space_intersection_point ) +{ + const float INTERSECTION_THRESHOLD = 0.001f; + + // +x face. + if ( object_space_intersection_point.x > 0.5f - INTERSECTION_THRESHOLD && + object_space_intersection_point.x < 0.5f + INTERSECTION_THRESHOLD ) { + return glm::vec3( 1.0f, 0.0f, 0.0f ); + } + // -x face. + else if ( object_space_intersection_point.x > -0.5f - INTERSECTION_THRESHOLD && + object_space_intersection_point.x < -0.5f + INTERSECTION_THRESHOLD ) { + return glm::vec3( -1.0f, 0.0f, 0.0f ); + } + // +y face. + else if ( object_space_intersection_point.y > 0.5f - INTERSECTION_THRESHOLD && + object_space_intersection_point.y < 0.5f + INTERSECTION_THRESHOLD ) { + return glm::vec3( 0.0f, 1.0f, 0.0f ); + } + // -y face. + else if ( object_space_intersection_point.y > -0.5f - INTERSECTION_THRESHOLD && + object_space_intersection_point.y < -0.5f + INTERSECTION_THRESHOLD ) { + return glm::vec3( 0.0f, -1.0f, 0.0f ); + } + // +z face. + else if ( object_space_intersection_point.z > 0.5f - INTERSECTION_THRESHOLD && + object_space_intersection_point.z < 0.5f + INTERSECTION_THRESHOLD ) { + return glm::vec3( 0.0f, 0.0f, 1.0f ); + } + // -z face. + else { + return glm::vec3( 0.0f, 0.0f, -1.0f ); + } +} + + // LOOK: Here's an intersection test example from a sphere. Now you just need to figure out cube and, optionally, triangle. // Sphere intersection test, return -1 if no intersection, otherwise, distance to intersection -__host__ __device__ float sphereIntersectionTest(staticGeom sphere, ray r, glm::vec3& intersectionPoint, glm::vec3& normal){ - - float radius = .5; - - glm::vec3 ro = multiplyMV(sphere.inverseTransform, glm::vec4(r.origin,1.0f)); - glm::vec3 rd = glm::normalize(multiplyMV(sphere.inverseTransform, glm::vec4(r.direction,0.0f))); +__host__ +__device__ +float sphereIntersectionTest( staticGeom sphere, + ray r, + glm::vec3& intersectionPoint, + glm::vec3& normal ) +{ + float radius = 0.5f; + + // Transform ray origin and direction to object space. + glm::vec3 ro = multiplyMV( sphere.inverseTransform, glm::vec4( r.origin, 1.0f ) ); + glm::vec3 rd = glm::normalize( multiplyMV( sphere.inverseTransform, glm::vec4( r.direction, 0.0f ) ) ); - ray rt; rt.origin = ro; rt.direction = rd; + ray rt; + rt.origin = ro; + rt.direction = rd; - float vDotDirection = glm::dot(rt.origin, rt.direction); - float radicand = vDotDirection * vDotDirection - (glm::dot(rt.origin, rt.origin) - pow(radius, 2)); - if (radicand < 0){ - return -1; - } + float vDotDirection = glm::dot( rt.origin, rt.direction ); + float radicand = vDotDirection * vDotDirection - ( glm::dot( rt.origin, rt.origin ) - pow( radius, 2.0f ) ); + if ( radicand < 0 ) { + return -1.0f; + } - float squareRoot = sqrt(radicand); - float firstTerm = -vDotDirection; - float t1 = firstTerm + squareRoot; - float t2 = firstTerm - squareRoot; + float squareRoot = sqrt( radicand ); + float firstTerm = -vDotDirection; + float t1 = firstTerm + squareRoot; + float t2 = firstTerm - squareRoot; - float t = 0; - if (t1 < 0 && t2 < 0) { - return -1; - } else if (t1 > 0 && t2 > 0) { - t = min(t1, t2); - } else { - t = max(t1, t2); - } - - glm::vec3 realIntersectionPoint = multiplyMV(sphere.transform, glm::vec4(getPointOnRay(rt, t), 1.0)); - glm::vec3 realOrigin = multiplyMV(sphere.transform, glm::vec4(0,0,0,1)); - - intersectionPoint = realIntersectionPoint; - normal = glm::normalize(realIntersectionPoint - realOrigin); + float t = 0.0f; + if ( t1 < 0.0f && t2 < 0.0f ) { + return -1.0f; + } else if ( t1 > 0.0f && t2 > 0.0f ) { + t = min( t1, t2 ); + } else { + t = max( t1, t2 ); + } + + glm::vec3 realIntersectionPoint = multiplyMV( sphere.transform, glm::vec4( getPointOnRay( rt, t ), 1.0f ) ); + glm::vec3 realOrigin = multiplyMV( sphere.transform, glm::vec4( 0.0f, 0.0f, 0.0f, 1.0f ) ); + + intersectionPoint = realIntersectionPoint; + normal = glm::normalize( realIntersectionPoint - realOrigin ); - return glm::length(r.origin - realIntersectionPoint); + return glm::length( r.origin - realIntersectionPoint ); } + // Returns x,y,z half-dimensions of tightest bounding box __host__ __device__ glm::vec3 getRadiuses(staticGeom geom){ glm::vec3 origin = multiplyMV(geom.transform, glm::vec4(0,0,0,1)); @@ -128,11 +301,12 @@ __host__ __device__ glm::vec3 getRadiuses(staticGeom geom){ return glm::vec3(xradius, yradius, zradius); } + // LOOK: Example for generating a random point on an object using thrust. // Generates a random point on a given cube __host__ __device__ glm::vec3 getRandomPointOnCube(staticGeom cube, float randomSeed){ - thrust::default_random_engine rng(hash(randomSeed)); + thrust::default_random_engine rng(simpleHash(randomSeed)); thrust::uniform_real_distribution u01(0,1); thrust::uniform_real_distribution u02(-0.5,0.5); @@ -174,13 +348,128 @@ __host__ __device__ glm::vec3 getRandomPointOnCube(staticGeom cube, float random } + // TODO: IMPLEMENT THIS FUNCTION // Generates a random point on a given sphere -__host__ __device__ glm::vec3 getRandomPointOnSphere(staticGeom sphere, float randomSeed){ +__host__ +__device__ +glm::vec3 getRandomPointOnSphere( staticGeom sphere, + float random_seed ) +{ + // Generate two random numbers u and v (0, 1) using thrust. + + thrust::default_random_engine rng( simpleHash( random_seed ) ); + thrust::uniform_real_distribution u01( 0.0f, 1.0f ); + + // Compute spherical coordinates theta and phi. + // theta = 2 * pi * u + // phi = inverse_cos( 2 * v - 1 ) + + float theta = 2.0f * PI * ( float )u01( rng ); + float phi = acos( 2.0f * ( float )u01( rng ) - 1.0f ); - return glm::vec3(0,0,0); + // Convert spherical coordinates (theta and phi) into cartesian coordinates (x, y, z). + // x = radius * sin( phi ) * cos( theta ) + // y = radius * sin( phi ) * sin( theta ) + // z = radius * cos( phi ) + + float x = 0.5f * sin( phi ) * cos( theta ); + float y = 0.5f * sin( phi ) * sin( theta ); + float z = 0.5f * cos( phi ); + + return multiplyMV( sphere.transform, glm::vec4( x, y, z, 1.0f) ); } -#endif + +__host__ +__device__ +glm::vec2 computeSphereUVCoordinates( staticGeom sphere, + glm::vec3 intersection_point ) +{ + // Thanks to https://www.cs.unc.edu/~rademach/xroads-RT/RTarticle.html#texturemap + + // Convert intersection point from world-space to object-space. + glm::vec3 Vp = glm::normalize( multiplyMV( sphere.inverseTransform, glm::vec4( intersection_point, 1.0f ) ) ); + + // Define unit vector pointing from the center of the sphere to the "north pole". + glm::vec3 Vn( 0.0f, 1.0f, 0.0f ); + + // Define unit vector pointing from the center of the sphere to any point on the equator. + glm::vec3 Ve( 1.0f, 0.0f, 0.0f ); + + // Compute angle between Vp and Vn, or the latitude. + float phi = acos( -glm::dot( Vn, Vp ) ); + + // Compute v, or the vertical image-space location [0, 1]. + float v = phi / PI; + + // Compute angle between Vp and Ve, or the longitude. + //float theta = acos( glm::dot( Vp, Ve ) / sin( phi ) ) / TWO_PI; + + float theta; + if ( sin( phi ) > -EPSILON && sin( phi ) < EPSILON ) { + theta = 0.0f; + } + else { + theta = acos( glm::dot( Vp, Ve ) / sin( phi ) ) / TWO_PI; + } + + // Compute u, or the horizontal image location [0, 1]: + float u; + if ( glm::dot( glm::cross( Vn, Ve ), Vp ) > 0.0f ) { + //u = theta; + u = 1.0f - theta; + } + else { + //u = 1.0f - theta; + u = theta; + } + + return glm::vec2( u, v ); +} + + +__host__ +__device__ +glm::vec2 computeCubeUVCoordinates( staticGeom cube, + glm::vec3 intersection_point ) +{ + // Convert intersection point from world-space to object-space. + glm::vec3 p = glm::normalize( multiplyMV( cube.inverseTransform, glm::vec4( intersection_point, 1.0f ) ) ); + + // Compute face ( 1 => x, 2 => y, 3 => z ). + int face = ( abs( p.x ) > abs( p.y ) && abs( p.x ) > abs( p.z ) ) ? 1 : ( abs( p.y ) > abs( p.z ) ) ? 2 : 3; + face = ( p[face - 1] < 0 ) ? ( -1.0f * face ) : face; + + float u, v; + + if ( face == 1 ) { // +x + u = ( ( p.z / abs( p.x ) ) + 1.0f ) / 2.0f; + v = ( ( p.y / abs( p.x ) ) + 1.0f ) / 2.0f; + } + else if ( face == -1 ) { // -x + u = ( ( -p.z / abs( p.x ) ) + 1.0f ) / 2.0f; + v = ( ( p.y / abs( p.x ) ) + 1.0f ) / 2.0f; + } + else if ( face == 2 ) { // +y + u = ( ( p.x / abs( p.y ) ) + 1.0f ) / 2.0f; + v = ( ( p.z / abs( p.y ) ) + 1.0f ) / 2.0f; + } + else if ( face == -2 ) { // -y + u = ( ( p.x / abs( p.y ) ) + 1.0f ) / 2.0f; + v = ( ( -p.z / abs( p.y ) ) + 1.0f ) / 2.0f; + } + else if ( face == 3 ) { // +z + u = ( ( p.x / abs( p.z ) ) + 1.0f ) / 2.0f; + v = ( ( p.y / abs( p.z ) ) + 1.0f ) / 2.0f; + } + else if ( face == -3 ) { // -z + u = ( ( -p.x / abs( p.z ) ) + 1.0f ) / 2.0f; + v = ( ( p.y / abs( p.z ) ) + 1.0f ) / 2.0f; + } + + return glm::vec2( u, v ); +} +#endif \ No newline at end of file diff --git a/src/main.cpp b/src/main.cpp index b002500..fa9a613 100755 --- a/src/main.cpp +++ b/src/main.cpp @@ -8,6 +8,21 @@ #include "main.h" #define GLEW_STATIC +#include "SimpleTimer.h" // Timer. + + +//////////////////////////////////////////////////// +// Constants +//////////////////////////////////////////////////// + +const std::string SCENE_FILE_NAME = "sampleScene_02.txt"; +const std::string OUTPUT_IMAGE_PATH = "C:\\Users\\Danny\\Documents\\_projects\\cis565\\Project3-Pathtracer\\renders\\"; +//const std::string TEXTURE_PATH = "C:\\Users\\Danny\\Documents\\_projects\\cis565\\Project3-Pathtracer\\data\\textures\\world_map.bmp"; +//const std::string SCENE_PATH = "C:\\Users\\Danny\\Documents\\_projects\\cis565\\Project3-Pathtracer\\data\\scenes\\sampleScene_02.txt"; + +SimpleTimer timer; // Timer. + + //------------------------------- //-------------MAIN-------------- //------------------------------- @@ -26,8 +41,9 @@ int main(int argc, char** argv){ istringstream liness(argv[i]); getline(liness, header, '='); getline(liness, data, '='); if(strcmp(header.c_str(), "scene")==0){ - renderScene = new scene(data); - loadedScene = true; + std::string full_scene_path = data + SCENE_FILE_NAME; // Danny was here. + renderScene = new scene(full_scene_path); + loadedScene = true; }else if(strcmp(header.c_str(), "frame")==0){ targetFrame = atoi(data.c_str()); singleFrameMode = true; @@ -36,6 +52,7 @@ int main(int argc, char** argv){ if(!loadedScene){ cout << "Error: scene file needed!" << endl; + std::cin.ignore(); return 0; } @@ -60,11 +77,14 @@ int main(int argc, char** argv){ } void mainLoop() { + + timer.start(); // Timer. + while(!glfwWindowShouldClose(window)){ glfwPollEvents(); runCuda(); - string title = "CIS565 Render | " + utilityCore::convertIntToString(iterations) + " Iterations"; + string title = "Danny Rerucha's Pathtracer | " + utilityCore::convertIntToString(iterations) + " Iterations"; glfwSetWindowTitle(window, title.c_str()); glBindBuffer(GL_PIXEL_UNPACK_BUFFER, pbo); @@ -94,9 +114,10 @@ void runCuda(){ iterations++; cudaGLMapBufferObject((void**)&dptr, pbo); - // pack geom and material arrays + // pack geom, material, and texture arrays geom* geoms = new geom[renderScene->objects.size()]; material* materials = new material[renderScene->materials.size()]; + simpleTexture *texture_images = new simpleTexture[renderScene->textures.size()]; for (int i=0; i < renderScene->objects.size(); i++) { geoms[i] = renderScene->objects[i]; @@ -104,14 +125,30 @@ void runCuda(){ for (int i=0; i < renderScene->materials.size(); i++) { materials[i] = renderScene->materials[i]; } - + for (int i=0; i < renderScene->textures.size(); i++) { + texture_images[i] = renderScene->textures[i]; + } + // execute the kernel - cudaRaytraceCore(dptr, renderCam, targetFrame, iterations, materials, renderScene->materials.size(), geoms, renderScene->objects.size() ); + cudaRaytraceCore( dptr, + renderCam, + targetFrame, + iterations, + materials, + renderScene->materials.size(), + geoms, + renderScene->objects.size(), + texture_images, + renderScene->textures.size() ); // unmap buffer object cudaGLUnmapBufferObject(pbo); } else { + float elapsed_time = timer.stop(); // Timer. + std::cout << "Elapsed time = " << elapsed_time << std::endl; // Timer. + std::cin.ignore(); // Timer. + if (!finishedRender) { // output image file image outputImage(renderCam->resolution.x, renderCam->resolution.y); @@ -119,7 +156,12 @@ void runCuda(){ for (int x=0; x < renderCam->resolution.x; x++) { for (int y=0; y < renderCam->resolution.y; y++) { int index = x + (y * renderCam->resolution.x); - outputImage.writePixelRGB(renderCam->resolution.x-1-x,y,renderCam->image[index]); + glm::vec3 pixel_color = renderCam->image[index]; + pixel_color.x /= iterations; + pixel_color.y /= iterations; + pixel_color.z /= iterations; + + outputImage.writePixelRGB(renderCam->resolution.x-1-x,y,pixel_color); } } @@ -133,8 +175,9 @@ void runCuda(){ stringstream out; out << targetFrame; s = out.str(); - utilityCore::replaceString(filename, ".bmp", "."+s+".bmp"); + utilityCore::replaceString(filename, ".bmp", "."+s+".bmp"); utilityCore::replaceString(filename, ".png", "."+s+".png"); + filename = OUTPUT_IMAGE_PATH + filename; // Danny was here. outputImage.saveImageRGB(filename); cout << "Saved frame " << s << " to " << filename << endl; finishedRender = true; @@ -170,7 +213,7 @@ bool init(int argc, char* argv[]) { width = 800; height = 800; - window = glfwCreateWindow(width, height, "CIS 565 Pathtracer", NULL, NULL); + window = glfwCreateWindow(width, height, "Danny Rerucha's Pathtracer", NULL, NULL); if (!window){ glfwTerminate(); return false; diff --git a/src/raytraceKernel.cu b/src/raytraceKernel.cu index 9c7bc7d..eab3ed9 100644 --- a/src/raytraceKernel.cu +++ b/src/raytraceKernel.cu @@ -9,6 +9,9 @@ #include #include +#include +#include + #include "sceneStructs.h" #include "glm/glm.hpp" #include "utilities.h" @@ -16,150 +19,721 @@ #include "intersections.h" #include "interactions.h" -void checkCUDAError(const char *msg) { - cudaError_t err = cudaGetLastError(); - if( cudaSuccess != err) { - fprintf(stderr, "Cuda error: %s: %s.\n", msg, cudaGetErrorString( err) ); - exit(EXIT_FAILURE); - } -} + +// Some new types. +enum MaterialType { + IDEAL_DIFFUSE, + PERFECT_SPUCULAR, + GLASS +}; + + +// Some forward declarations. +__host__ __device__ bool sceneIntersection( const ray &r, staticGeom *geoms, int num_geoms, float &t, int &id, glm::vec3 &intersection_point, glm::vec3 &intersection_normal ); +__host__ __device__ MaterialType determineMaterialType( material mat ); +__host__ __device__ glm::vec2 computePixelSubsampleLocation( glm::vec2 pixel_index, glm::vec2 image_resolution, int current_iteration ); + + +void checkCUDAError( const char *msg ) +{ + cudaError_t err = cudaGetLastError(); + if ( cudaSuccess != err) { + fprintf( stderr, "Cuda error: %s: %s.\n", msg, cudaGetErrorString( err ) ); + + // DEBUG. + std::cin.ignore(); + + exit( EXIT_FAILURE ); + } +} + // LOOK: This function demonstrates how to use thrust for random number generation on the GPU! // Function that generates static. -__host__ __device__ glm::vec3 generateRandomNumberFromThread(glm::vec2 resolution, float time, int x, int y){ - int index = x + (y * resolution.x); - - thrust::default_random_engine rng(hash(index*time)); - thrust::uniform_real_distribution u01(0,1); +__host__ +__device__ +glm::vec3 generateRandomNumberFromThread( glm::vec2 resolution, + float time, + int x, + int y ) +{ + int index = x + ( y * resolution.x ); + //float simple_hash_seed = ( float )( ( int )( index * time ) % INT_MAX ); + //float simple_hash_seed = ( float )( ( int )( index * time ) % 1000000 ); - return glm::vec3((float) u01(rng), (float) u01(rng), (float) u01(rng)); -} + thrust::default_random_engine rng( simpleHash( index * time ) ); + //thrust::default_random_engine rng( simpleHash( simple_hash_seed ) ); + thrust::uniform_real_distribution u01( 0, 1 ); -// TODO: IMPLEMENT THIS FUNCTION -// Function that does the initial raycast from the camera -__host__ __device__ ray raycastFromCameraKernel(glm::vec2 resolution, float time, int x, int y, glm::vec3 eye, glm::vec3 view, glm::vec3 up, glm::vec2 fov){ - ray r; - r.origin = glm::vec3(0,0,0); - r.direction = glm::vec3(0,0,-1); - return r; + return glm::vec3( ( float )u01( rng ), + ( float )u01( rng ), + ( float )u01( rng ) ); } -//Kernel that blacks out a given image buffer -__global__ void clearImage(glm::vec2 resolution, glm::vec3* image){ - int x = (blockIdx.x * blockDim.x) + threadIdx.x; - int y = (blockIdx.y * blockDim.y) + threadIdx.y; - int index = x + (y * resolution.x); - if(x<=resolution.x && y<=resolution.y){ - image[index] = glm::vec3(0,0,0); - } + +// Kernel that blacks out a given image buffer +__global__ +void clearImage( glm::vec2 resolution, + glm::vec3* image ) +{ + int x = ( blockIdx.x * blockDim.x ) + threadIdx.x; + int y = ( blockIdx.y * blockDim.y ) + threadIdx.y; + int index = x + ( y * resolution.x ); + if ( x <= resolution.x && y <= resolution.y ) { + image[index] = glm::vec3( 0, 0, 0 ); + } } -//Kernel that writes the image to the OpenGL PBO directly. -__global__ void sendImageToPBO(uchar4* PBOpos, glm::vec2 resolution, glm::vec3* image){ - - int x = (blockIdx.x * blockDim.x) + threadIdx.x; - int y = (blockIdx.y * blockDim.y) + threadIdx.y; - int index = x + (y * resolution.x); + +// Kernel that writes the image to the OpenGL PBO directly. +__global__ +void sendImageToPBO( uchar4* PBOpos, + glm::vec2 resolution, + glm::vec3* image, + int iterations_so_far ) +{ + int x = ( blockIdx.x * blockDim.x ) + threadIdx.x; + int y = ( blockIdx.y * blockDim.y ) + threadIdx.y; + int index = x + ( y * resolution.x ); - if(x<=resolution.x && y<=resolution.y){ + if ( x <= resolution.x && y <= resolution.y ) { + glm::vec3 color; + color.x = ( image[index].x / iterations_so_far ) * 255.0; + color.y = ( image[index].y / iterations_so_far ) * 255.0; + color.z = ( image[index].z / iterations_so_far ) * 255.0; + + if ( color.x > 255 ) { + color.x = 255; + } + if ( color.y > 255 ) { + color.y = 255; + } + if ( color.z > 255 ) { + color.z = 255; + } + + // Each thread writes one pixel location in the texture (textel) + PBOpos[index].w = 0; + PBOpos[index].x = color.x; + PBOpos[index].y = color.y; + PBOpos[index].z = color.z; + } +} - glm::vec3 color; - color.x = image[index].x*255.0; - color.y = image[index].y*255.0; - color.z = image[index].z*255.0; - if(color.x>255){ - color.x = 255; - } +/*********** CORE PATHTRACING ALGORITHMS ***********/ - if(color.y>255){ - color.y = 255; - } - if(color.z>255){ - color.z = 255; - } - - // Each thread writes one pixel location in the texture (textel) - PBOpos[index].w = 0; - PBOpos[index].x = color.x; - PBOpos[index].y = color.y; - PBOpos[index].z = color.z; - } +__host__ +__device__ +bool sceneIntersection( const ray &r, + staticGeom *geoms, + int num_geoms, + float &t, + int &id, + glm::vec3 &intersection_point, + glm::vec3 &intersection_normal ) +{ + t = FLT_MAX; + float temp_t = -1.0f; + glm::vec3 temp_intersection_point; + glm::vec3 temp_intersection_normal; + + // Find nearest intersection, if any. + for ( int i = 0; i < num_geoms; ++i ) { + if ( geoms[i].type == SPHERE ) { + temp_t = sphereIntersectionTest( geoms[i], + r, + temp_intersection_point, + temp_intersection_normal ); + } + else if ( geoms[i].type == CUBE ) { + temp_t = boxIntersectionTest( geoms[i], + r, + temp_intersection_point, + temp_intersection_normal ); + } + + // Update nearest intersection if closer intersection has been found. + if ( temp_t > 0.0f && temp_t < t ) { + t = temp_t; + intersection_point = temp_intersection_point; + intersection_normal = temp_intersection_normal; + //id = geoms[i].materialid; + id = i; + } + } + + return ( t < FLT_MAX ); +} + + +//__host__ +//__device__ +//bool lightIntersection( const ray &r, +// staticGeom *geoms, +// int num_geoms, +// int &obj_id ) +//{ +// float t = FLT_MAX; +// float temp_t = -1.0f; +// glm::vec3 temp_intersection_point; +// glm::vec3 temp_intersection_normal; +// +// // Find nearest intersection, if any. +// for ( int i = 0; i < num_geoms; ++i ) { +// if ( geoms[i].type == SPHERE ) { +// temp_t = sphereIntersectionTest( geoms[i], +// r, +// temp_intersection_point, +// temp_intersection_normal ); +// } +// else if ( geoms[i].type == CUBE ) { +// temp_t = boxIntersectionTest( geoms[i], +// r, +// temp_intersection_point, +// temp_intersection_normal ); +// } +// +// // Update nearest intersection if closer intersection has been found. +// if ( temp_t > 0.0f && temp_t < t ) { +// t = temp_t; +// obj_id = i; +// } +// } +// +// return ( t < FLT_MAX ); +//} + + +// Compute rays from camera through pixels and store in ray_pool. +__global__ +void raycastFromCameraKernel( ray *ray_pool, + glm::vec2 resolution, + glm::vec3 eyep, + glm::vec3 m, + glm::vec3 h, + glm::vec3 v, + int current_iteration ) +{ + int x = ( blockIdx.x * blockDim.x ) + threadIdx.x; + int y = ( blockIdx.y * blockDim.y ) + threadIdx.y; + int index = ( y * ( int )resolution.x ) + x; + + if ( index > ( resolution.x * resolution.y ) ) { + return; + } + + // TODO: Add a switch to turn anti-aliasing on/off. + glm::vec2 pixel_location = computePixelSubsampleLocation( glm::vec2( x, y ), resolution, current_iteration ); + + float sx = pixel_location.x / ( resolution.x - 1.0f ); + float sy = 1.0f - ( pixel_location.y / ( resolution.y - 1.0f ) ); + + //float sx = ( float )x / ( resolution.x - 1.0f ); + //float sy = 1.0f - ( ( float )y / ( resolution.y - 1.0f ) ); + + glm::vec3 image_point = m + ( ( 2.0f * sx - 1.0f ) * h ) + ( ( 2.0f * sy - 1.0f ) * v ); + + glm::vec3 dir = image_point - eyep; + + ray r; + r.origin = eyep; + r.direction = glm::normalize( dir ); + r.image_coords = glm::vec2( x, y ); + + ray_pool[index] = r; } -// TODO: IMPLEMENT THIS FUNCTION -// Core raytracer kernel -__global__ void raytraceRay(glm::vec2 resolution, float time, cameraData cam, int rayDepth, glm::vec3* colors, - staticGeom* geoms, int numberOfGeoms){ - int x = (blockIdx.x * blockDim.x) + threadIdx.x; - int y = (blockIdx.y * blockDim.y) + threadIdx.y; - int index = x + (y * resolution.x); +__host__ +__device__ +glm::vec2 computePixelSubsampleLocation( glm::vec2 pixel_index, + glm::vec2 image_resolution, + int current_iteration ) +{ + // TODO: Set pixel index (x, y) as center of pixel instead of bottom left corner. + // TODO: Get rid of loop b/c it's completely unnecessary and wastes resources. + + const int SUBSAMPLE_ROWS = 4; + const int SUBSAMPLE_COLS = 4; + + float row_height = 1.0f / SUBSAMPLE_ROWS; + float col_width = 1.0f / SUBSAMPLE_COLS; - if((x<=resolution.x && y<=resolution.y)){ + float left_bounds, upper_bounds; + float x_percentage, y_percentage, new_x, new_y; - colors[index] = generateRandomNumberFromThread(resolution, time, x, y); - } + int sub_pixel_num = current_iteration % ( SUBSAMPLE_ROWS * SUBSAMPLE_COLS ); + + // partition pixel into grid using perfect_square and randomly sample one ray through each grid space + upper_bounds = ( float )pixel_index.y; + for ( int y = 0; y < SUBSAMPLE_ROWS; ++y ) { + left_bounds = ( float )pixel_index.x; + for ( int x = 0; x < SUBSAMPLE_COLS; ++x ) { + if ( ( y * SUBSAMPLE_COLS ) + x == sub_pixel_num ) { + glm::vec3 rand = generateRandomNumberFromThread( image_resolution, current_iteration, pixel_index.x, pixel_index.y ); + x_percentage = rand.x; + y_percentage = rand.y; + + new_x = left_bounds + ( col_width * x_percentage ); + new_y = upper_bounds + ( row_height * y_percentage ); + + return glm::vec2( new_x, new_y ); + } + left_bounds += col_width; + } + upper_bounds += row_height; + } + return pixel_index; } -// TODO: FINISH THIS FUNCTION -// Wrapper for the __global__ call that sets up the kernel calls and does a ton of memory management -void cudaRaytraceCore(uchar4* PBOpos, camera* renderCam, int frame, int iterations, material* materials, int numberOfMaterials, geom* geoms, int numberOfGeoms){ - - int traceDepth = 1; //determines how many bounces the raytracer traces - // set up crucial magic - int tileSize = 8; - dim3 threadsPerBlock(tileSize, tileSize); - dim3 fullBlocksPerGrid((int)ceil(float(renderCam->resolution.x)/float(tileSize)), (int)ceil(float(renderCam->resolution.y)/float(tileSize))); +// Core raytracer kernel. +__global__ +void raytraceRay( ray *ray_pool, + int ray_pool_size, + glm::vec2 resolution, + float current_iteration, // Used solely for random number generation (I think). + cameraData cam, + int raytrace_depth, + glm::vec3 *image_buffer, + staticGeom *geoms, + int num_geoms, + material *materials, + int *texture_info, + glm::vec3 *texture_rgb ) +{ + int ray_pool_index = ( blockIdx.x * blockDim.x ) + threadIdx.x; + + if ( ray_pool_index > ray_pool_size ) { + return; + } + + ray r = ray_pool[ray_pool_index]; + int image_pixel_index = ( r.image_coords.y * ( int )resolution.x ) + r.image_coords.x; + + // Only necessary when stream compaction is turned off. + if ( !r.is_active ) { + return; + } + + // Nudge ray along it's direction to avoid intersecting with the surface it originates from. + r.origin += ( r.direction * 0.001f ); + + // Intersection testing. + float dist_to_intersection; + //int material_index; + int intersected_obj_id; + glm::vec3 intersection_point; + glm::vec3 intersection_normal; + bool ray_did_intersect_something = sceneIntersection( r, // Current ray. + geoms, // List of scene geometry. + num_geoms, // Number of pieces of geometry. + dist_to_intersection, // Reference to be filled. + intersected_obj_id, // Reference to be filled. + intersection_point, // Reference to be filled. + intersection_normal ); // Reference to be filled. + + + // Ray misses. Return background color. Kill ray. + if ( !ray_did_intersect_something ) { + image_buffer[image_pixel_index] += glm::vec3( 0.0f, 0.0f, 0.0f ); + r.is_active = false; + ray_pool[ray_pool_index] = r; + return; + } + + // Get intersected object and material of intersected object. + staticGeom obj = geoms[intersected_obj_id]; + material mat = materials[obj.materialid]; + + // Use Roussian Roulette to randomly kill the current ray. + // rgb is the color of the intersected material, and is used in the computations that follow. + glm::vec3 rgb = mat.color; + float p = ( rgb.x > rgb.y && rgb.x > rgb.z ) ? rgb.x : ( ( rgb.y > rgb.z ) ? rgb.y : rgb.z ); + if ( raytrace_depth > 5 ) { + glm::vec3 rand = generateRandomNumberFromThread( resolution, ( current_iteration * raytrace_depth ), r.image_coords.x, r.image_coords.y ); + if ( rand.x < p ) { + rgb = rgb * ( 1.0f / p ); + } + else { + image_buffer[image_pixel_index] += ( r.color * rgb * mat.emittance ); + r.is_active = false; + ray_pool[ray_pool_index] = r; + return; + } + } + + + // Intersected object has a texture image. + //glm::vec3 foobar = rgb; + if ( obj.texture_id != -1 ) { + + glm::vec2 uv( 0.0f, 0.0f ); + + if ( obj.type == SPHERE ) { + uv = computeSphereUVCoordinates( obj, intersection_point ); + } + else if ( obj.type == CUBE ) { + uv = computeCubeUVCoordinates( obj, intersection_point ); + } + + // Extract texture info. + int texture_info_base_index = obj.texture_id * 3; + int texture_rgb_starting_index = texture_info[texture_info_base_index]; + int texture_width = texture_info[texture_info_base_index + 1]; + int texture_height = texture_info[texture_info_base_index + 2]; + + // Compute pixel coordinates within single texture. + int pixel_coord_x = ( int )( uv.x * texture_width ); + int pixel_coord_y = ( int )( uv.y * texture_height ); + + if ( !( pixel_coord_x < 0 || pixel_coord_y < 0 || pixel_coord_x >= texture_width || pixel_coord_y >= texture_height ) ) { + // Compute texture RGB index. + int pixel_linear_index = ( pixel_coord_y * texture_width ) + pixel_coord_x; + int texture_rgb_index = texture_rgb_starting_index + pixel_linear_index; + + if ( texture_rgb_index < texture_rgb_starting_index || texture_rgb_index >= ( texture_rgb_starting_index + texture_width * texture_height ) ) { + rgb = glm::vec3( 1.0f, 0.078f, 0.577f ); // Hot pink to visualize an error. + } + else { + rgb = texture_rgb[texture_rgb_index]; + } + } + else { + rgb = glm::vec3( 1.0f, 0.078f, 0.577f ); // Hot pink to visualize an error. + } + + } + + // Ray hits light source. Add acculumated color contribution of ray. Kill ray. + if ( mat.emittance > 0.0f ) { + image_buffer[image_pixel_index] += ( r.color * rgb * mat.emittance ); + r.is_active = false; + ray_pool[ray_pool_index] = r; + return; + } + + MaterialType mat_type = determineMaterialType( mat ); + if ( mat_type == IDEAL_DIFFUSE ) { + glm::vec3 rand = generateRandomNumberFromThread( resolution, + current_iteration * raytrace_depth, + r.image_coords.x, + r.image_coords.y ); + r.direction = calculateRandomDirectionInHemisphere( intersection_normal, + rand.x, + rand.y ); + + //// Compute direct illumination contribution. + //glm::vec3 e( 0.0f, 0.0f, 0.0f ); + //for ( int i = 0; i < num_geoms; ++i ) { + // const staticGeom &s = geoms[i]; + // material l_mat = materials[s.materialid]; + + // // Skip geometry that isn't a light source. + // if ( l_mat.emittance < 0.0f ) { + // continue; + // } + + // glm::vec3 light_dir; + // if ( s.type == SPHERE ) { + // light_dir = glm::normalize( getRandomPointOnSphere( s, ( current_iteration * raytrace_depth ) ) - intersection_point ); + // } + // else if ( s.type == CUBE ) { + // light_dir = glm::normalize( getRandomPointOnCube( s, ( current_iteration * raytrace_depth ) ) - intersection_point ); + // } + // else if ( s.type == MESH ) { + // // TODO. + // } + // else { + // // ERROR: Unrecognized geometry type. + // } + + // ray light_ray; + // light_ray.direction = light_dir; + // light_ray.origin = intersection_point; + + // // Intersection testing. + // int light_id; + // bool did_hit_light = lightIntersection( light_ray, geoms, num_geoms, light_id ); + + // glm::vec3 intersection_normal_oriented = ( glm::dot( intersection_normal, r.direction ) < 0.0f ) ? intersection_normal : ( -1.0f * intersection_normal ); + + // if ( did_hit_light && light_id == i ) { + // e += f * ( l_mat.emittance * glm::dot( light_dir, intersection_normal_oriented ) ); + // } + //} + + r.color = r.color * rgb; // Only add color contributions of this ray if it makes contact with a light source. + r.origin = intersection_point; // Set origin point for next ray. + ray_pool[ray_pool_index] = r; // Update ray in ray pool. + return; + } + else if ( mat_type == PERFECT_SPUCULAR ) { + r.direction = calculateReflectionDirection( intersection_normal, r.direction ); // Mirror surface contributes no color. + r.origin = intersection_point; // Set origin point for next ray. + ray_pool[ray_pool_index] = r; // Update ray in ray pool. + return; + } + else if ( mat_type == GLASS ) { + glm::vec3 intersection_normal_oriented = ( glm::dot( intersection_normal, r.direction ) < 0.0f ) ? intersection_normal : ( -1.0f * intersection_normal ); + bool ray_is_entering = ( glm::dot( intersection_normal, intersection_normal_oriented ) > 0.0f ); + + const float IOR_GLASS = 1.5f; + float ior_incident = ray_is_entering ? 1.0f : IOR_GLASS; + float ior_transmitted = ray_is_entering ? IOR_GLASS : 1.0f; + + glm::vec3 refl_dir = calculateReflectionDirection( intersection_normal, r.direction ); + glm::vec3 trans_dir = calculateTransmissionDirection( intersection_normal, r.direction, ior_incident, ior_transmitted ); + + Fresnel fres = calculateFresnel( intersection_normal, r.direction, ior_incident, ior_transmitted, refl_dir, trans_dir ); + + glm::vec3 rand = generateRandomNumberFromThread( resolution, ( current_iteration * raytrace_depth ), r.image_coords.x, r.image_coords.y ); + + if ( rand.x < fres.reflectionCoefficient ) { + r.direction = refl_dir; + r.origin = intersection_point; + ray_pool[ray_pool_index] = r; + return; + } + else { + r.direction = trans_dir; + r.origin = intersection_point; + ray_pool[ray_pool_index] = r; + return; + } + } +} + + +__host__ +__device__ +MaterialType determineMaterialType( material mat ) +{ + if ( mat.hasRefractive ) { + return GLASS; + } + else if ( mat.hasReflective ) { + return PERFECT_SPUCULAR; + } + else { + return IDEAL_DIFFUSE; + } +} + + +// thrust predicate to cull inactive rays from ray pool. +struct RayIsInactive +{ + __host__ + __device__ + bool operator()( const ray &r ) + { + return !r.is_active; + } +}; + + +// Wrapper that sets up kernel calls and handles memory management. +// Handles one pathtrace iteration. Called many times to produce a rendered image. +void cudaRaytraceCore( uchar4 *pbo_pos, + camera *render_cam, + int frame, + int current_iteration, + material *materials, + int num_materials, + geom *geoms, + int num_geoms, + simpleTexture *textures, + int num_textures ) +{ + // Tune these for performance. + int depth = 10; + int camera_raycast_tile_size = 8; + int raytrace_tile_size = 128; + + // Setup crucial magic. + dim3 threads_per_block( camera_raycast_tile_size, + camera_raycast_tile_size ); + dim3 full_blocks_per_grid( ( int )ceil( ( float )render_cam->resolution.x / ( float )camera_raycast_tile_size ), + ( int )ceil( ( float )render_cam->resolution.y / ( float )camera_raycast_tile_size ) ); - // send image to GPU - glm::vec3* cudaimage = NULL; - cudaMalloc((void**)&cudaimage, (int)renderCam->resolution.x*(int)renderCam->resolution.y*sizeof(glm::vec3)); - cudaMemcpy( cudaimage, renderCam->image, (int)renderCam->resolution.x*(int)renderCam->resolution.y*sizeof(glm::vec3), cudaMemcpyHostToDevice); + // Send image to GPU. + glm::vec3 *cuda_image = NULL; + float size_image = ( int )render_cam->resolution.x * ( int )render_cam->resolution.y * sizeof( glm::vec3 ); + cudaMalloc( ( void** )&cuda_image, + size_image ); + cudaMemcpy( cuda_image, + render_cam->image, + size_image, + cudaMemcpyHostToDevice ); - // package geometry and materials and sent to GPU - staticGeom* geomList = new staticGeom[numberOfGeoms]; - for(int i=0; iresolution; - cam.position = renderCam->positions[frame]; - cam.view = renderCam->views[frame]; - cam.up = renderCam->ups[frame]; - cam.fov = renderCam->fov; + // Send materials to GPU. + material *cuda_materials = NULL; + float size_material_list = num_materials * sizeof( material ); + cudaMalloc( ( void** )&cuda_materials, + size_material_list ); + cudaMemcpy( cuda_materials, + materials, + size_material_list, + cudaMemcpyHostToDevice ); - // kernel launches - raytraceRay<<>>(renderCam->resolution, (float)iterations, cam, traceDepth, cudaimage, cudageoms, numberOfGeoms); - sendImageToPBO<<>>(PBOpos, renderCam->resolution, cudaimage); + // TODO: The below computations for textures only need to be done once. - // retrieve image from GPU - cudaMemcpy( renderCam->image, cudaimage, (int)renderCam->resolution.x*(int)renderCam->resolution.y*sizeof(glm::vec3), cudaMemcpyDeviceToHost); + // Create int *texture_info. + int *texture_info = new int[num_textures * 3]; // Starting index, width, height. + int starting_index = 0; + for ( int i = 0; i < num_textures; ++i ) { + int ti = i * 3; + texture_info[ti] = starting_index; + texture_info[ti + 1] = textures[i].width; + texture_info[ti + 2] = textures[i].height; + starting_index += ( textures[i].width * textures[i].height ); + } - // free up stuff, or else we'll leak memory like a madman - cudaFree( cudaimage ); - cudaFree( cudageoms ); - delete geomList; + // Create glm::vec3 *texture_rgb. + int num_texture_pixels = starting_index; + glm::vec3 *texture_rgb = new glm::vec3[num_texture_pixels]; + int rgb_index = 0; + for ( int i = 0; i < num_textures; ++i ) { + for ( int y = 0; y < textures[i].height; ++y ) { + for ( int x = 0; x < textures[i].width; ++x ) { + texture_rgb[rgb_index] = ( textures[i].rgb[( y * textures[i].width ) + x] ); + ++rgb_index; + } + } + } - // make certain the kernel has completed - cudaThreadSynchronize(); + // Send texture_info and texture_rgb to CUDA. + int *cuda_texture_info = NULL; + int size_texture_info = num_textures * 3 * sizeof( int ); + cudaMalloc( ( void** )&cuda_texture_info, + size_texture_info ); + cudaMemcpy( cuda_texture_info, + texture_info, + size_texture_info, + cudaMemcpyHostToDevice ); - checkCUDAError("Kernel failed!"); -} + glm::vec3 *cuda_texture_rgb = NULL; + int size_texture_rgb = num_texture_pixels * sizeof( glm::vec3 ); + cudaMalloc( ( void** )&cuda_texture_rgb, + size_texture_rgb ); + cudaMemcpy( cuda_texture_rgb, + texture_rgb, + size_texture_rgb, + cudaMemcpyHostToDevice ); + + // Package up camera. + cameraData cam; + cam.resolution = render_cam->resolution; + cam.position = render_cam->positions[frame]; + cam.view = render_cam->views[frame]; + cam.up = render_cam->ups[frame]; + cam.fov = render_cam->fov; + + // Variables to compute rays originating from render camera. + glm::vec3 a = glm::cross( cam.view, cam.up ); + glm::vec3 m = cam.position + cam.view; // Midpoint of frame buffer. + glm::vec3 h = ( a * glm::length( cam.view ) * ( float )tan( cam.fov.x * ( PI / 180.0f ) ) ) / glm::length( a ); // Horizontal NDC value. + glm::vec3 v = glm::vec3( 0.0f, cam.resolution.y * glm::length( h ) / cam.resolution.x, 0.0f ); // Vertical NDC value. + + // Allocate device memory for ray pool. + ray *cuda_ray_pool = NULL; + int num_rays = ( int )( render_cam->resolution.x * render_cam->resolution.y ); + cudaMalloc( ( void** )&cuda_ray_pool, + num_rays * sizeof( ray ) ); + + // Initialize ray pool with rays originating at the render camera directed through each pixel in the image buffer. + raycastFromCameraKernel<<< full_blocks_per_grid, threads_per_block >>>( cuda_ray_pool, + render_cam->resolution, + cam.position, + m, + h, + v, + ( float )current_iteration ); + + // Launch raytraceRay kernel once per raytrace depth. + for ( int i = 0; i < depth; ++i ) { + dim3 threads_per_raytrace_block( raytrace_tile_size ); + dim3 blocks_per_raytrace_grid( ( int )ceil( ( float )num_rays / ( float )raytrace_tile_size ) ); + + // Launch raytraceRay kernel. + raytraceRay<<< blocks_per_raytrace_grid, threads_per_raytrace_block >>>( cuda_ray_pool, + num_rays, + render_cam->resolution, + ( float )current_iteration, + cam, + ( i + 1 ), + cuda_image, + cuda_geoms, + num_geoms, + cuda_materials, + cuda_texture_info, + cuda_texture_rgb ); + + // Note: Stream compaction is slow. + thrust::device_ptr ray_pool_device_ptr( cuda_ray_pool ); + thrust::device_ptr culled_ray_pool_device_ptr = thrust::remove_if( ray_pool_device_ptr, + ray_pool_device_ptr + num_rays, + RayIsInactive() ); + + // Compute number of active rays in ray pool. + num_rays = culled_ray_pool_device_ptr.get() - ray_pool_device_ptr.get(); + } + + // Launch sendImageToPBO kernel. + sendImageToPBO<<< full_blocks_per_grid, threads_per_block >>>( pbo_pos, + render_cam->resolution, + cuda_image, + current_iteration ); + + // Retrieve image from GPU. + cudaMemcpy( render_cam->image, + cuda_image, + size_image, + cudaMemcpyDeviceToHost ); + + // Cleanup. + cudaFree( cuda_image ); + cudaFree( cuda_geoms ); + cudaFree( cuda_materials ); + cudaFree( cuda_ray_pool ); + cudaFree( cuda_texture_info ); + cudaFree( cuda_texture_rgb ); + delete[] geom_list; + delete[] texture_info; + delete[] texture_rgb; + + // Make certain the kernel has completed. + //cudaThreadSynchronize(); // Deprecated. + cudaDeviceSynchronize(); + + checkCUDAError( "Kernel failed!" ); +} \ No newline at end of file diff --git a/src/raytraceKernel.h b/src/raytraceKernel.h index 984e89f..627220f 100755 --- a/src/raytraceKernel.h +++ b/src/raytraceKernel.h @@ -13,7 +13,17 @@ #include #include #include "sceneStructs.h" +#include "image.h" -void cudaRaytraceCore(uchar4* pos, camera* renderCam, int frame, int iterations, material* materials, int numberOfMaterials, geom* geoms, int numberOfGeoms); +void cudaRaytraceCore( uchar4 *pbo_pos, + camera *render_cam, + int frame, + int current_iteration, + material *materials, + int num_materials, + geom* geoms, + int num_geoms, + simpleTexture *textures, + int num_textures ); #endif diff --git a/src/scene.cpp b/src/scene.cpp index 4cbe216..d132091 100644 --- a/src/scene.cpp +++ b/src/scene.cpp @@ -29,6 +29,10 @@ scene::scene(string filename){ loadCamera(); cout << " " << endl; } + else if ( strcmp( tokens[0].c_str(), "TEXTURE" ) == 0 ) { + loadTextures( tokens[1] ); + std::cout << " " << std::endl; + } } } } @@ -73,11 +77,21 @@ int scene::loadObject(string objectid){ //link material utilityCore::safeGetline(fp_in,line); - if(!line.empty() && fp_in.good()){ + if(!line.empty() && fp_in.good()) { vector tokens = utilityCore::tokenizeString(line); newObject.materialid = atoi(tokens[1].c_str()); cout << "Connecting Object " << objectid << " to Material " << newObject.materialid << "..." << endl; - } + } + + // Attach textures. + utilityCore::safeGetline( fp_in,line ); + if( !line.empty() && fp_in.good() ) { + vector tokens = utilityCore::tokenizeString( line ); + if ( strcmp( tokens[0].c_str(), "texture" ) == 0 ) { + newObject.texture_id = atoi( tokens[1].c_str() ); + cout << "Loading texture " << newObject.texture_id << " for object " << objectid << "..." << endl; + } + } //load frames int frameCount = 0; @@ -263,3 +277,62 @@ int scene::loadMaterial(string materialid){ return 1; } } + + +const std::string TEXTURE_PATH = "C:\\Users\\Danny\\Documents\\_projects\\cis565\\Project3-Pathtracer\\data\\textures\\"; + +int scene::loadTextures( std::string texture_id ) +{ + int id = atoi( texture_id.c_str() ); + + if ( id != textures.size() ) { + std::cout << "ERROR: TEXTURE ID does not match expected number of textures." << std::endl; + return -1; + } + else{ + cout << "Loading texture " << id << "..." << endl; + + std::string line; + utilityCore::safeGetline( fp_in, line ); + std::vector tokens = utilityCore::tokenizeString( line ); + + std::string texture_filename = ""; + if ( strcmp( tokens[0].c_str(), "FILE" ) == 0 ) { + texture_filename = tokens[1]; + } + + std::string texture_full_path = TEXTURE_PATH + texture_filename; + + int i; + FILE *f = fopen( texture_full_path.c_str(), "rb" ); + unsigned char info[54]; + fread( info, sizeof( unsigned char ), 54, f ); // read the 54-byte header + + // extract image height and width from header + int width = *( int* )&info[18]; + int height = *( int* )&info[22]; + + int size = 3 * width * height; + unsigned char* data = new unsigned char[size]; // allocate 3 bytes per pixel + fread( data, sizeof( unsigned char ), size, f ); // read the rest of the data at once + fclose( f ); + + simpleTexture texture_image; + texture_image.width = width; + texture_image.height = height; + + texture_image.rgb = new glm::vec3[width * height]; + + for ( int y = 0; y < height; ++y ) { + for ( int x = 0; x < width; ++x ) { + int linear_index = ( y * width ) + x; + int bmp_buffer_loc = linear_index * 3; + glm::vec3 rgb( data[bmp_buffer_loc + 1] / 255.0f, data[bmp_buffer_loc] / 255.0f, data[bmp_buffer_loc + 2] / 255.0f ); + texture_image.rgb[linear_index] = rgb; + } + } + + textures.push_back( texture_image ); + return 1; + } +} \ No newline at end of file diff --git a/src/scene.h b/src/scene.h index ea3ddaa..c9607c5 100644 --- a/src/scene.h +++ b/src/scene.h @@ -22,6 +22,7 @@ class scene{ int loadMaterial(string materialid); int loadObject(string objectid); int loadCamera(); + int loadTextures( std::string texture_id ); public: scene(string filename); ~scene(); @@ -29,6 +30,8 @@ class scene{ vector objects; vector materials; camera renderCam; + + std::vector textures; }; #endif diff --git a/src/sceneStructs.h b/src/sceneStructs.h index 5e0c853..7af75dc 100644 --- a/src/sceneStructs.h +++ b/src/sceneStructs.h @@ -14,8 +14,21 @@ enum GEOMTYPE{ SPHERE, CUBE, MESH }; struct ray { + __host__ + __device__ + ray(): + is_active( true ), + color( glm::vec3( 1.0f, 1.0f, 1.0f ) ), + index_of_refraction( 1.0f ) + {} + glm::vec3 origin; glm::vec3 direction; + + glm::vec3 color; + glm::vec2 image_coords; + float index_of_refraction; + bool is_active; }; struct geom { @@ -27,6 +40,7 @@ struct geom { glm::vec3* scales; cudaMat4* transforms; cudaMat4* inverseTransforms; + int texture_id; }; struct staticGeom { @@ -37,6 +51,7 @@ struct staticGeom { glm::vec3 scale; cudaMat4 transform; cudaMat4 inverseTransform; + int texture_id; }; struct cameraData { @@ -73,4 +88,10 @@ struct material{ float emittance; }; +struct simpleTexture { + int width; + int height; + glm::vec3 *rgb; +}; + #endif //CUDASTRUCTS_H diff --git a/src/utilities.cpp b/src/utilities.cpp index a8e5d90..02c47a3 100755 --- a/src/utilities.cpp +++ b/src/utilities.cpp @@ -4,7 +4,7 @@ // File: utilities.cpp // A collection/kitchen sink of generally useful functions -#define GLM_FORCE_RADIANS +//#define GLM_FORCE_RADIANS #include #include @@ -137,6 +137,8 @@ std::istream& utilityCore::safeGetline(std::istream& is, std::string& t) } } } + + //----------------------------- //-------GLM Printers---------- //----------------------------- diff --git a/src/utilities.h b/src/utilities.h index f51598f..1e84fda 100644 --- a/src/utilities.h +++ b/src/utilities.h @@ -16,6 +16,7 @@ #include #include #include "cudaMat4.h" +#include "image.h" #define PI 3.1415926535897932384626422832795028841971 #define TWO_PI 6.2831853071795864769252867665590057683943 diff --git a/windows/Project3-Pathtracer/Project3-Pathtracer/Project3-Pathtracer.vcxproj b/windows/Project3-Pathtracer/Project3-Pathtracer/Project3-Pathtracer.vcxproj index c45dd79..97296ee 100644 --- a/windows/Project3-Pathtracer/Project3-Pathtracer/Project3-Pathtracer.vcxproj +++ b/windows/Project3-Pathtracer/Project3-Pathtracer/Project3-Pathtracer.vcxproj @@ -28,7 +28,7 @@ - + @@ -77,6 +77,7 @@ + @@ -86,6 +87,7 @@ + @@ -95,6 +97,6 @@ - + \ No newline at end of file diff --git a/windows/Project3-Pathtracer/Project3-Pathtracer/Project3-Pathtracer.vcxproj.filters b/windows/Project3-Pathtracer/Project3-Pathtracer/Project3-Pathtracer.vcxproj.filters index 584fd19..66306b6 100644 --- a/windows/Project3-Pathtracer/Project3-Pathtracer/Project3-Pathtracer.vcxproj.filters +++ b/windows/Project3-Pathtracer/Project3-Pathtracer/Project3-Pathtracer.vcxproj.filters @@ -48,6 +48,9 @@ Header Files + + Header Files + @@ -71,6 +74,9 @@ Source Files\stb_image + + Source Files +