diff --git a/PROJ1_WIN/565Raytracer/565Raytracer.vcxproj b/PROJ1_WIN/565Raytracer/565Raytracer.vcxproj index 201fd23..3ab0b55 100755 --- a/PROJ1_WIN/565Raytracer/565Raytracer.vcxproj +++ b/PROJ1_WIN/565Raytracer/565Raytracer.vcxproj @@ -75,7 +75,7 @@ - + @@ -91,10 +91,10 @@ - true + false - true + false false @@ -198,8 +198,23 @@ C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.5\include;C:/ProgramData/NVIDIA Corporation/CUDA Samples/v5.5/common/inc;../shared/glew/includes;../shared/freeglut/includes + + + C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.5\include;C:\ProgramData\NVIDIA Corporation\CUDA Samples\v5.5\common\inc;../shared/glew/include;../shared/freeglut/include;%(AdditionalIncludeDirectories) + + + + + false + + + + + C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.5\include;C:\ProgramData\NVIDIA Corporation\CUDA Samples\v5.5\common\inc;../shared/glew/includes;../shared/freeglut/includes + + - + \ No newline at end of file diff --git a/PROJ1_WIN/565Raytracer/README_images/001_flat_shading.bmp b/PROJ1_WIN/565Raytracer/README_images/001_flat_shading.bmp new file mode 100644 index 0000000..577f8a1 Binary files /dev/null and b/PROJ1_WIN/565Raytracer/README_images/001_flat_shading.bmp differ diff --git a/PROJ1_WIN/565Raytracer/README_images/002_diffuse_illumination.bmp b/PROJ1_WIN/565Raytracer/README_images/002_diffuse_illumination.bmp new file mode 100644 index 0000000..f8b9aa5 Binary files /dev/null and b/PROJ1_WIN/565Raytracer/README_images/002_diffuse_illumination.bmp differ diff --git a/PROJ1_WIN/565Raytracer/README_images/003_diffuse_illumination_with_hard_shadows.bmp b/PROJ1_WIN/565Raytracer/README_images/003_diffuse_illumination_with_hard_shadows.bmp new file mode 100644 index 0000000..436f6e6 Binary files /dev/null and b/PROJ1_WIN/565Raytracer/README_images/003_diffuse_illumination_with_hard_shadows.bmp differ diff --git a/PROJ1_WIN/565Raytracer/README_images/004_diffuse_illumination_with_soft_shadows.bmp b/PROJ1_WIN/565Raytracer/README_images/004_diffuse_illumination_with_soft_shadows.bmp new file mode 100644 index 0000000..5067408 Binary files /dev/null and b/PROJ1_WIN/565Raytracer/README_images/004_diffuse_illumination_with_soft_shadows.bmp differ diff --git a/PROJ1_WIN/565Raytracer/README_images/005_phong_illumination_with_soft_shadows_and_reflections.bmp b/PROJ1_WIN/565Raytracer/README_images/005_phong_illumination_with_soft_shadows_and_reflections.bmp new file mode 100644 index 0000000..4d775c3 Binary files /dev/null and b/PROJ1_WIN/565Raytracer/README_images/005_phong_illumination_with_soft_shadows_and_reflections.bmp differ diff --git a/PROJ1_WIN/565Raytracer/README_images/005_phong_illumination_with_soft_shadows_and_reflections_and_supersampled_antialiasing.bmp b/PROJ1_WIN/565Raytracer/README_images/005_phong_illumination_with_soft_shadows_and_reflections_and_supersampled_antialiasing.bmp new file mode 100644 index 0000000..0243abc Binary files /dev/null and b/PROJ1_WIN/565Raytracer/README_images/005_phong_illumination_with_soft_shadows_and_reflections_and_supersampled_antialiasing.bmp differ diff --git a/PROJ1_WIN/565Raytracer/README_images/006_final.bmp b/PROJ1_WIN/565Raytracer/README_images/006_final.bmp new file mode 100644 index 0000000..da6b196 Binary files /dev/null and b/PROJ1_WIN/565Raytracer/README_images/006_final.bmp differ diff --git a/PROJ1_WIN/565Raytracer/README_images/Thumbs.db b/PROJ1_WIN/565Raytracer/README_images/Thumbs.db new file mode 100644 index 0000000..1fc4c57 Binary files /dev/null and b/PROJ1_WIN/565Raytracer/README_images/Thumbs.db differ diff --git a/PROJ1_WIN/565Raytracer/README_images/block_data_chart.bmp b/PROJ1_WIN/565Raytracer/README_images/block_data_chart.bmp new file mode 100644 index 0000000..2e409a4 Binary files /dev/null and b/PROJ1_WIN/565Raytracer/README_images/block_data_chart.bmp differ diff --git a/PROJ1_WIN/565Raytracer/README_images/block_data_graph.bmp b/PROJ1_WIN/565Raytracer/README_images/block_data_graph.bmp new file mode 100644 index 0000000..555f6b1 Binary files /dev/null and b/PROJ1_WIN/565Raytracer/README_images/block_data_graph.bmp differ diff --git a/PROJ1_WIN/565Raytracer/README_images/weird.bmp b/PROJ1_WIN/565Raytracer/README_images/weird.bmp new file mode 100644 index 0000000..79a70fe Binary files /dev/null and b/PROJ1_WIN/565Raytracer/README_images/weird.bmp differ diff --git a/README.md b/README.md index 6bef2b9..7c9ff70 100644 --- a/README.md +++ b/README.md @@ -1,129 +1,234 @@ ------------------------------------------------------------------------------- -CIS565: Project 1: CUDA Raytracer +CIS565 Project 1: CUDA Raytracer ------------------------------------------------------------------------------- -Fall 2013 +Ricky Arietta Fall 2013 ------------------------------------------------------------------------------- -Due Thursday, 09/19/2013 ------------------------------------------------------------------------------- +This project is a highly parallel version of the traditional ray tracing +algorithm on the implemented on the GPU, augmented from a template provided by +Karl Yi and Liam Boone. It works by casting out virtual rays through an image plane +from a user-defined camera and assigning pixel values based on intersections +with the defined geometry and lights. The final implementation, when built, +is capable of rendering high-quality images composed of physically realistic +area light(s) and geometry primitives (spheres, cubes). An example of a final +render can be seen immediately below this description. As you can see, the final +implementation accounts for Phong Illumination, soft shadows from area lights, +recursive reflections of light within the scene, and supersampled antialiasing. The +implementation of each feature is described briefly below with a rendered image +demonstrating the development of the code. Finally, there is some performance +analysis included regarding the size and number of blocks requested on the GPU +during runtime. + +(A brief tour of the base code and a description of the scene file format +used during implementation are also included at the end of this file, adapted +from the project description provided by Patrick Cozzi and Liam Boone.) + +![Final Rendered Image](https://raw.github.com/rarietta/Project1-RayTracer/master/PROJ1_WIN/565Raytracer/README_images/006_final.bmp) + ------------------------------------------------------------------------------- -NOTE: +Initial Ray Casting From Camera & Geometry Primitive Intersection ------------------------------------------------------------------------------- -This project requires an NVIDIA graphics card with CUDA capability! Any card -after the Geforce 8xxx series will work. If you do not have an NVIDIA graphics -card in the machine you are working on, feel free to use any machine in the SIG -Lab or in Moore100 labs. All machines in the SIG Lab and Moore100 are equipped -with CUDA capable NVIDIA graphics cards. If this too proves to be a problem, -please contact Patrick or Liam as soon as possible. + +In order to render images using ray tracing, we needed to define the set of rays +being sent out from the camera. The camera, defined by the user input file, +includes a resolution, field of view angle, viewing direction, up direction, etc. +Using this information, I was able to define an arbitrary image plane some distance +from the camera, orthogonal to the viewing direction vector. Then, using the +up direction and a computed third orthogonal "right" vector, I was able to +define a grid on the image plane with the same resolution as the desired image. +Then, from the camera, and given an x-index and y-index for the pixel in question, +I could compute a single ray from the camera position to the center of the corresponding +pixel in the image plane (adapted for supersampled antialiasing; see below). These +rays served as the initial rays for tracing paths through the scene. + +Additionally, when following these rays, I needed to be able to determine any +geometry intersections along the ray direction vector, since these intersections +define the illumination color value returned to the image pixel. Thus, in addition +to the provided sphereIntersectionTest(), I created a boxIntersectionTest() for +computing ray collisions with cube primitives. This tested for an intersection with +the plane defined by each face, tested it against the size of the cube to see if +the intersection was bounded by the edges, and returned the minimum-distance intersection +from this set. + +A very early render is seen below, proving the successful implementation of ray-casting +and primitive intersection into a multi-colored Cornell box with 3 spheres of. This +render only displays raycasting and intersection, and thus no lighting model was implemented. +All the surfaces are rendered according to their flat diffuse RGB color value. + +![Flat Shading](https://raw.github.com/rarietta/Project1-RayTracer/master/PROJ1_WIN/565Raytracer/README_images/001_flat_shading.bmp) ------------------------------------------------------------------------------- -INTRODUCTION: +Addition of Diffuse Lambertian Shading with Point Lighting ------------------------------------------------------------------------------- -In this project, you will implement a CUDA based raytracer capable of -generating raytraced rendered images extremely quickly. For those of you who -have taken CIS460/560, building a raytracer should not be anything new to you -from a conceptual point of 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). - -The ultimate purpose of this project is to serve as the foundation for your -next project: a full CUDA based global illumination pathtracer. Raytracing can -be thought of as a way to generate an isolated version of the direct light -contribution in a global illumination scenario. - -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 raytracer, 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. + +The earliest step in creating a lighting model within the ray tracer was to +implement Lambertian Diffuse lighting for the geometry surfaces. This model of +lighting follows Lambert's equations, which take into account the direction +from the intersection to the light and the normal of the geometry at the +intersection point. Thus, when the normal and the light are in the same direction, +the luminance value is greatest. When they point in opposite directions, the +light has no contribution to the pixel luminance at that point. + +An example of this diffuse lighting model is seen below. You can see that, compared +to the flat shading model in which the white sources all blended together in the +image plane, the geometry now has some semblance of volume and boundaries can be +determined by the eye. + +![Flat Shading](https://raw.github.com/rarietta/Project1-RayTracer/master/PROJ1_WIN/565Raytracer/README_images/002_diffuse_illumination.bmp) ------------------------------------------------------------------------------- -CONTENTS: +Addition of Hard Shadows with Point Lighting ------------------------------------------------------------------------------- -The Project1 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. -* scenes/ contains an example scene description file. -* renders/ contains an example render of the given example scene file. -* PROJ1_WIN/ contains a Windows Visual Studio 2010 project and all dependencies - needed for building and running on Windows 7. -* PROJ1_OSX/ contains a OSX makefile, run script, and all dependencies needed - for building and running on Mac OSX 10.8. -* PROJ1_NIX/ contains a Linux makefile for building and running on Ubuntu - 12.04 LTS. Note that you will need to set the following environment - variables: - - - PATH=$PATH:/usr/local/cuda-5.5/bin - - LD_LIBRARY_PATH=/usr/local/cuda-5.5/lib64:/lib - - you may set these any way that you like. I added them to my .bashrc - - -The Windows and OSX versions of the project build and run exactly the same way -as in Project0. + +This addition was one of the first steps in generating shadows within the scene and +imitating the behavior of light. I initially treated each light source as a simple +point light. Thus, for every intersection with the scene from a camera ray, a secondary +shadow ray or shadow feeler was cast in the direction of the constant point light +source. If this ray intersected any other geometry before reaching the light source, +it was considered in shadow and the luminance value of the pixel was set to zero +(black). If it did not intersect any additional geometry, then the light source +had a direct path to the point and the luminance was calculated in the normal fashion. + +An example of this point light/hard shadow model is included below. As you can see +in the next section, this model was quickly adapted to account for geometric area lights +and soft shadows, but this image is a good indicator of the progression of the code. + +![Flat Shading](https://raw.github.com/rarietta/Project1-RayTracer/master/PROJ1_WIN/565Raytracer/README_images/003_diffuse_illumination_with_hard_shadows.bmp) ------------------------------------------------------------------------------- -REQUIREMENTS: +Addition of Soft Shadows with Area Lighting ------------------------------------------------------------------------------- -In this project, you are given code for: - -* Loading, reading, and storing the TAKUAscene 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 -* Phong lighting for one point light source -* Diffuse lambertian surfaces -* Raytraced shadows -* Cube intersection testing -* Sphere surface point sampling - -You are also required to implement at least 2 of the following features: - -* Specular reflection -* Soft shadows and area lights -* Texture mapping -* Bump mapping -* Depth of field -* Supersampled antialiasing -* Refraction, i.e. glass -* OBJ Mesh loading and renderin -* Interactive camera + +To adapt the illumination model to account for soft shadows and area lights, +rather than unrealistic hard shadows and single point lights, I simply sampled +the light source randomly over multiple iterations. For each iteration, a +shadow ray was cast to a randomly generated point on the cube (or sphere) light +source geometry (using the required functions defined in the project assignment). +A shadow ray was cast from the intersection point to this random light source, +and a contibution of the luminance would be averaged into the pixel value for that +iteration. If the shadow ray intersected geometry and thus could not see the light +point, no luminance contribution was added during that iteration. Thus, over time, +the shadow term described above became a fractional value, rather than a simple zero +or one. + +A rendered image with soft shadows is seen below. You can see that this result +is much more physically based and realistic, accounting for the physical light +source casting light rays in multiple directions towards/around the geometry. + +![Flat Shading](https://raw.github.com/rarietta/Project1-RayTracer/master/PROJ1_WIN/565Raytracer/README_images/004_diffuse_illumination_with_soft_shadows.bmp) + +------------------------------------------------------------------------------- +Addition of Phong Illumination and Reflective Ray Tracing +------------------------------------------------------------------------------- + +In addition to the Diffuse Lambertian lighting model, Phong illumination was +implemented. Thus, the Lambertian term is summed with a specular term, defined +by the specular RGB value of the material, the angle of specular reflection, +and the specular exponent. This implementation also accounted for linear +attenuation of the light sources. Thus, the light's influence on the luminance +value was scaled inversely with the distance from the light source. *(Since no +linear attenuation coefficient was provided in the scene file format used in +this project, I performed this operation with an arbitrarily chosen constant +that was equal for all light sources and seemed to produce an aesthetically +pleasing final result).* + +Additionally, recursive specular ray tracing was implemented, which mimics +the behavior of real, specular, reflective surfaces. While this is traditionally +done recursively in ray tracing, CUDA does not support recursive function calls. +Thus, a color value was stored and for each intersection, if the material was +reflective, a secondary ray was traced out and the secondary luminance was added +to the first color value, scaled by the specular coefficient of the material. +This was implemented in a for-loop, performed a number of times equal to the +specified traceDepth of the program. + +The below image is an example of both specular Phong model (including linear +attenuation of the light sources) and specular reflection with a traceDepth of 2. + +![Flat Shading](https://raw.github.com/rarietta/Project1-RayTracer/master/PROJ1_WIN/565Raytracer/README_images/005_phong_illumination_with_soft_shadows_and_reflections.bmp) + +------------------------------------------------------------------------------- +Addition of Supersampled AntiAliasing +------------------------------------------------------------------------------- + +With the existing code base described up to this point, it was easy to implement +antialiasing by supersampling the pixel values. Instead of casting the same ray +through the center of the pixel with each iteration, the direction of the ray +within the bounds of the pixel were determined randomly in each iteration, and +the computed intersection illumination values were averaged over the entire series +of iterations (much like the implementation of area lighting). + +Compare the following image to the image in the previous section. All input +and scene data was identical between the two, except this version included +supersampling of the pixels. You can see how smooth the edges are on the spheres +and cubes in this version. While there are clear "jaggies" in the above version, +the below version has none and even corrects for tricky edge intersection +cases in the corners of the box. + +![Flat Shading](https://raw.github.com/rarietta/Project1-RayTracer/master/PROJ1_WIN/565Raytracer/README_images/005_phong_illumination_with_soft_shadows_and_reflections_and_supersampled_antialiasing.bmp) + +This implementation is clearly superior. The random sampling of the pixels does, +however, make it impossible to store the first level intersections in a spatial +data structure for easy lookup in each iteration, which may be a technique for +accelerating the ray tracing runtime. However, this is a tradeoff for vastly +superior image quality, and I believe an important addition to the code. + +------------------------------------------------------------------------------- +PERFORMANCE EVALUATION +------------------------------------------------------------------------------- + +To analyze the performance of the program on the GPU hardware, I decided to run +timing tests on the renders with varying block sizes/counts. To do this, I altered +the tile size within the rendered image, increasing or decreasing the block size +and count when invoking the kernel. + +*Caveat*: Unfortunately, the hardware to which I had access provided certain limitations +to my testing. The following block sizes were the only ones that the hardware in +the Moore 100 Lab could handle. I intended to test the code with more tile sizes +in [1,20], but they all caused the hardware to crash. Furthermore, I will +add that my program seemed to run at a generally slower rate than the demonstrations +provided in class, and I believe this was an artifact of the hardware as well. Regardless, +the following data shows certain trends in hardware utilization, it just lacks the +completeness I would have desired. + +![Flat Shading](https://raw.github.com/rarietta/Project1-RayTracer/master/PROJ1_WIN/565Raytracer/README_images/block_data_chart.bmp) + +As you can see, a tile size of 8 when rendering these images seemed to be the most efficient. +Increasing the tile size to 10 slowed down the iteration renders dramatically, indicating +an inefficient utilization of GPU memory and processing power. Interestingly enough, +doubling the size of the tile, while increasing the runtime, did not nearly have as much +of a negative impact on the runtime as an increase in tile size from 8 to 10. + +We can also see that the runtime per iteration increases with an increase in the +traceDepth of the rays, or how many times the rays are recursively followed through +a scene when calculating reflective contributions. This is to be expected, since more +instructions and memory accesses are required for deeper paths. But we can see that +the increase is not linear, having the greatest jump from traceDepth=1 (no reflection) +to traceDepth=2. + +The above data is visualized in a graph below, plotting the average render time per +iteration for each tileSize for each traceDepth: + +![Flat Shading](https://raw.github.com/rarietta/Project1-RayTracer/master/PROJ1_WIN/565Raytracer/README_images/block_data_graph.bmp) + +------------------------------------------------------------------------------- +Runtime Video +------------------------------------------------------------------------------- + +Unfortunately, since I was working in Moore 100, I was unable to download or +utilize and screen capture video software for producing runtime videos. ------------------------------------------------------------------------------- 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. +The main files of interest in this prooject, which handle the ray-tracing +algorithm and image generation, are the following: * 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 raytracing - logic should be implemented in this CUDA kernel. raytraceRay() should + * cudaRaytraceCore() handles kernel launches and memory management + * raycastFromCameraKernel() handles camera raycasting. + * raytraceRay() is the core raytracing 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. @@ -135,9 +240,7 @@ reference are marked with the comment LOOK. 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. + This function should work in the same way as getRandomPointOnCube(). * interactions.h contains functions for ray-object interactions that define how rays behave upon hitting materials and objects. You will need to complete: @@ -145,31 +248,9 @@ reference are marked with the comment LOOK. 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. + lights, cameras, and animation frames are stored in the renderer. ------------------------------------------------------------------------------- TAKUAscene FORMAT: @@ -247,87 +328,3 @@ 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. - -------------------------------------------------------------------------------- -README -------------------------------------------------------------------------------- -All students must replace or augment the contents of this Readme.md in a clear -manner with the following: - -* A brief description of the project and the specific features you implemented. -* At least one screenshot of your project running. -* A 30 second or longer video of your project running. To create the video you - can use http://www.microsoft.com/expression/products/Encoder4_Overview.aspx -* A performance evaluation (described in detail below). - -------------------------------------------------------------------------------- -PERFORMANCE EVALUATION -------------------------------------------------------------------------------- -The performance evaluation is where you will investigate how to make your CUDA -programs more efficient using the skills you've learned in class. You must have -perform at least one experiment on your code to investigate the positive or -negative effects on performance. - -One such experiment would be to investigate the performance increase involved -with adding a spatial data-structure to your scene data. - -Another idea could be looking at the change in timing between various block -sizes. - -A good metric to track would be number of rays per second, or frames per -second, or number of objects displayable at 60fps. - -We encourage you to get creative with your tweaks. Consider places in your code -that could be considered bottlenecks and try to improve them. - -Each student should provide no more than a one page summary of their -optimizations along with tables and or graphs to visually explain and -performance differences. - -------------------------------------------------------------------------------- -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 Liam, - liamboone+cis565@gmail.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 -------------------------------------------------------------------------------- -As with the previous project, you should fork this project and work inside of -your fork. Upon completion, commit your finished project back to your fork, and -make a pull request to the master repository. You should include a README.md -file in the root directory detailing the following - -* A brief description of the project and specific features you implemented -* At least one screenshot of your project running, and at least one screenshot - of the final rendered output of your raytracer -* A link to a video of your raytracer running. -* Instructions for building and running your project if they differ from the - base code -* A performance writeup as detailed above. -* A list of all third-party code used. -* This Readme file, augmented or replaced as described above in the README section. diff --git a/scenes/sampleScene.txt b/scenes/sampleScene.txt index 877706f..3a1d7c4 100755 --- a/scenes/sampleScene.txt +++ b/scenes/sampleScene.txt @@ -34,11 +34,11 @@ ABSCOEFF 0 0 0 RSCTCOEFF 0 EMITTANCE 0 -MATERIAL 3 //red glossy -RGB .63 .06 .04 -SPECEX 0 -SPECRGB 1 1 1 -REFL 0 +MATERIAL 3 //yellow glossy +RGB 1 1 0 +SPECEX 0.99999 +SPECRGB .6 .6 .6 +REFL 1 REFR 0 REFRIOR 2 SCATTER 0 @@ -48,9 +48,9 @@ EMITTANCE 0 MATERIAL 4 //white glossy RGB 1 1 1 -SPECEX 0 +SPECEX 0.9 SPECRGB 1 1 1 -REFL 0 +REFL 1 REFR 0 REFRIOR 2 SCATTER 0 @@ -70,11 +70,11 @@ ABSCOEFF .02 5.1 5.7 RSCTCOEFF 13 EMITTANCE 0 -MATERIAL 6 //green glossy -RGB .15 .48 .09 -SPECEX 0 +MATERIAL 6 //purple glossy +RGB .5 0 .5 +SPECEX 0.999 SPECRGB 1 1 1 -REFL 0 +REFL 1 REFR 0 REFRIOR 2.6 SCATTER 0 @@ -108,8 +108,8 @@ EMITTANCE 15 CAMERA RES 800 800 -FOVY 25 -ITERATIONS 5000 +FOVY 25 25 +ITERATIONS 500 FILE test.bmp frame 0 EYE 0 4.5 12 @@ -166,7 +166,7 @@ SCALE 3 3 3 OBJECT 6 sphere -material 3 +material 6 frame 0 TRANS 2 5 2 ROTAT 0 180 0 @@ -174,7 +174,7 @@ SCALE 2.5 2.5 2.5 OBJECT 7 sphere -material 6 +material 3 frame 0 TRANS -2 5 -2 ROTAT 0 180 0 diff --git a/src/interactions.h b/src/interactions.h index 8c3f5f6..9054afd 100755 --- a/src/interactions.h +++ b/src/interactions.h @@ -26,6 +26,8 @@ __host__ __device__ glm::vec3 calculateTransmissionDirection(glm::vec3 normal, g __host__ __device__ glm::vec3 calculateReflectionDirection(glm::vec3 normal, glm::vec3 incident); __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); +__host__ __device__ glm::vec3 computePhongTotal(ray& r, glm::vec3 intersection_point, glm::vec3 intersection_normal, material intersection_mtl, staticGeom* lights, int numberOfLights, staticGeom* geoms, int numberOfGeoms, material* materials, float time); +__host__ __device__ float computeShadowCoefficient(glm::vec3 intersection_point, staticGeom light, staticGeom* geoms, int numberOfGeoms, float time); //TODO (OPTIONAL): IMPLEMENT THIS FUNCTION __host__ __device__ glm::vec3 calculateTransmission(glm::vec3 absorptionCoefficient, float distance) { @@ -45,8 +47,12 @@ __host__ __device__ glm::vec3 calculateTransmissionDirection(glm::vec3 normal, g //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); + float IdotN = glm::dot(-incident,normal); + glm::vec3 I; + if (IdotN < 0.0f) { I = incident; } + else { I = -incident; } + glm::vec3 R = glm::normalize(2*IdotN*normal - I); + return R; } //TODO (OPTIONAL): IMPLEMENT THIS FUNCTION @@ -67,7 +73,8 @@ __host__ __device__ glm::vec3 calculateRandomDirectionInHemisphere(glm::vec3 nor float over = sqrt(1 - up * up); // sin(theta) float around = xi2 * TWO_PI; - //Find a direction that is not the normal based off of whether or not the normal's components are all equal to sqrt(1/3) or whether or not at least one component is less than sqrt(1/3). Learned this trick from Peter Kutz. + //Find a direction that is not the normal based off of whether or not the normal's components are all equal to sqrt(1/3) + //or whether or not at least one component is less than sqrt(1/3). Learned this trick from Peter Kutz. glm::vec3 directionNotNormal; if (abs(normal.x) < SQRT_OF_ONE_THIRD) { @@ -90,7 +97,118 @@ __host__ __device__ glm::vec3 calculateRandomDirectionInHemisphere(glm::vec3 nor //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); + + float z = xi1; + float theta = xi2 * TWO_PI; + + float r = sqrt(1-z*z); + float x = r*cos(theta); + float y = r*sin(theta); + + return glm::vec3(x,y,z); +} + +__host__ __device__ glm::vec3 computePhongTotal(ray& r, glm::vec3 intersection_point, glm::vec3 intersection_normal, + material intersection_mtl, staticGeom* lights, int numberOfLights, + staticGeom* geoms, int numberOfGeoms, material* materials, float time){ + glm::vec3 rgb(0.0, 0.0, 0.0); + + float n = intersection_mtl.specularExponent; + glm::vec3 ks = glm::vec3(intersection_mtl.specularColor); + glm::vec3 kd = glm::vec3(intersection_mtl.color); + glm::vec3 N = glm::vec3(intersection_normal); + glm::vec3 V = glm::vec3(-r.direction); + + glm::vec3 I, R, L; + + for (int i = 0; i < numberOfLights; i++) { + staticGeom light = lights[i]; + + float shadow_coefficient = computeShadowCoefficient(intersection_point, light, geoms, numberOfGeoms, time); + if (shadow_coefficient > 0.0) + { + // get point on light source + glm::vec3 light_point = getRandomPointOnGeom(light,time); + + // direction from intersection point to light source + L = glm::normalize(light_point - intersection_point); + float LdotN = glm::dot(L,N); + //if (LdotN<0) LdotN=0.0f; if (LdotN>1) LdotN=1.0f; + + // direction of perfect specular reflection + R = glm::normalize(2*LdotN*N - L); + float RdotV = glm::dot(R,V); + if (RdotV<0) RdotV=0.0f; if (RdotV>1) RdotV=1.0f; + + // light material + material light_mtl = materials[light.materialid]; + // light intensity + I = light_mtl.color * min((float)1.0, materials[light.materialid].emittance); + + // specular phong term + glm::vec3 specular; + if (n == 0.0f) { specular = glm::vec3(0.0f,0.0f,0.0f); } + else { specular = ks*pow(RdotV, n)*0.5f; } + if (specular.x > 1.0f) { specular.x = 1.0f; } else if (specular.x < 0.0f) { specular.x = 0.0f; } + if (specular.y > 1.0f) { specular.y = 1.0f; } else if (specular.y < 0.0f) { specular.y = 0.0f; } + if (specular.z > 1.0f) { specular.z = 1.0f; } else if (specular.z < 0.0f) { specular.z = 0.0f; } + specular *= glm::vec3(0.1f, 0.1f, 0.1f); //scale back specular component, rendering too bright + + // diffuse phong term + glm::vec3 diffuse = kd*LdotN; + + // phong reflectance model + rgb += (diffuse + specular) * shadow_coefficient * I; + rgb += intersection_mtl.emittance; + } + } + if (rgb.x > 1.0f) { rgb.x = 1.0f; } else if (rgb.x < 0.0f) { rgb.x = 0.0f; } + if (rgb.y > 1.0f) { rgb.y = 1.0f; } else if (rgb.y < 0.0f) { rgb.y = 0.0f; } + if (rgb.z > 1.0f) { rgb.z = 1.0f; } else if (rgb.z < 0.0f) { rgb.z = 0.0f; } + return rgb; +} + +__host__ __device__ float computeShadowCoefficient(glm::vec3 intersection_point, staticGeom light, + staticGeom* geoms, int numberOfGeoms, float randomSeed) { + + //glm::vec3 ro = glm::vec3(intersection_point); + //glm::vec3 rd = multiplyMV(light.transform, glm::vec4(0,0,0,1.0f)) - intersection_point; + //ray rt; rt.origin = ro; rt.direction = rd; + + thrust::default_random_engine rng(hash(randomSeed)); + thrust::uniform_real_distribution u01(-0.5,0.5); + glm::vec3 random_light_point = getRandomPointOnGeom(light,randomSeed); + glm::vec3 ro = glm::vec3(intersection_point); + glm::vec3 rd = random_light_point - intersection_point; + ray rt; rt.origin = ro; rt.direction = rd; + + glm::vec3 shadow_point; + glm::vec3 shadow_normal; + float dist_to_light = geomIntersectionTest(light, rt, shadow_point, shadow_normal); + + for (int i = 0; i < numberOfGeoms; i++) { + + // return values for intersection tests + float t = -1.0; + + // current geometry object + staticGeom geom = geoms[i]; + + // if the geometry is equivalent to the light source, skip it + if (geom.objectid == light.objectid) + continue; + + // test for intersections with sphere/box + t = geomIntersectionTest(geom, rt, shadow_point, shadow_normal); + + // see if geometry was intersected before the light source + float error = 1e-3; + if ((t > error) && (t < dist_to_light)) + return (0.0f); + } + //account for linear attenuation of light source (unless geom is light source) + if (dist_to_light > 1e-4) { return 1.0f/(0.25f*dist_to_light); } + else { return 1.0f; } } //TODO (PARTIALLY OPTIONAL): IMPLEMENT THIS FUNCTION diff --git a/src/intersections.h b/src/intersections.h index daefe95..b6350d6 100755 --- a/src/intersections.h +++ b/src/intersections.h @@ -19,7 +19,9 @@ __host__ __device__ glm::vec3 getSignOfRay(ray r); __host__ __device__ glm::vec3 getInverseDirectionOfRay(ray r); __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 getRandomPointOnGeom(staticGeom geom, float randomSeed); __host__ __device__ glm::vec3 getRandomPointOnCube(staticGeom cube, float randomSeed); +__host__ __device__ glm::vec3 getRandomPointOnSphere(staticGeom sphere, float randomSeed); //Handy dandy little hashing function that provides seeds for random number generation __host__ __device__ unsigned int hash(unsigned int a){ @@ -68,11 +70,72 @@ __host__ __device__ glm::vec3 getSignOfRay(ray r){ return glm::vec3((int)(inv_direction.x < 0), (int)(inv_direction.y < 0), (int)(inv_direction.z < 0)); } +//Cube intersection test, return -1 if no intersection, otherwise, distance to intersection +__host__ __device__ float geomIntersectionTest(staticGeom geom, ray r, glm::vec3& intersectionPoint, glm::vec3& normal){ + if (geom.type == CUBE) return boxIntersectionTest(geom, r, intersectionPoint, normal); + else if (geom.type == SPHERE) return sphereIntersectionTest(geom, r, intersectionPoint, normal); + return (float)-1.0; +} + //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){ + + 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; + + glm::vec3 faceNormals[6]; + glm::vec3 faceCenters[6]; + faceNormals[0] = glm::vec3(0,0,-1); faceCenters[0] = glm::vec3(0,0,-0.5); + faceNormals[1] = glm::vec3(0,0,-1); faceCenters[1] = glm::vec3(0,0, 0.5); + faceNormals[2] = glm::vec3(0,-1,0); faceCenters[2] = glm::vec3(0,-0.5,0); + faceNormals[3] = glm::vec3(0, 1,0); faceCenters[3] = glm::vec3(0, 0.5,0); + faceNormals[4] = glm::vec3(-1,0,0); faceCenters[4] = glm::vec3(-0.5,0,0); + faceNormals[5] = glm::vec3( 1,0,0); faceCenters[5] = glm::vec3( 0.5,0,0); + + // closest discovered intersection + float min_t = -1.0; + int min_i = 6; + + // find intersection of ray with each plane of the box + for (unsigned int i = 0; i < 6; i++) { + glm::vec3 normal = faceNormals[i]; + glm::vec3 center = faceCenters[i]; - return -1; + float t = glm::dot((center - rt.origin), normal) / glm::dot(rt.direction, normal); + + // continue if intersection is behind camera + if (t <= 0) + continue; + + // if t is greater than the closest found intersection, skip it + if ((min_t > 0.0) && (t >= min_t)) + continue; + + // check to see if the point found is within + // the edges defined by the face + glm::vec3 P = getPointOnRay(rt,t); + float error = 0.75e-3; + if ((P.x >= (-0.5 - error)) && (P.x <= (0.5 + error)) && + (P.y >= (-0.5 - error)) && (P.y <= (0.5 + error)) && + (P.z >= (-0.5 - error)) && (P.z <= (0.5 + error))) + min_t = t; + min_i = i; + } + + if (min_t < 0) + return (float) -1.0; + + else { + glm::vec3 realIntersectionPoint = multiplyMV(box.transform, glm::vec4(getPointOnRay(rt, min_t), 1.0)); + glm::vec3 realNormal = glm::normalize(multiplyMV(box.transform, glm::vec4(faceNormals[min_i],0.0f))); + intersectionPoint = realIntersectionPoint; + normal = realNormal; + + return glm::length(r.origin - realIntersectionPoint); + } } //LOOK: Here's an intersection test example from a sphere. Now you just need to figure out cube and, optionally, triangle. @@ -89,7 +152,7 @@ __host__ __device__ float sphereIntersectionTest(staticGeom sphere, ray r, glm:: 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; + return (float) -1.0; } float squareRoot = sqrt(radicand); @@ -127,6 +190,12 @@ __host__ __device__ glm::vec3 getRadiuses(staticGeom geom){ return glm::vec3(xradius, yradius, zradius); } +__host__ __device__ glm::vec3 getRandomPointOnGeom(staticGeom geom, float randomSeed){ + if (geom.type == SPHERE) { return getRandomPointOnSphere(geom, randomSeed); } + else if (geom.type == CUBE) { return getRandomPointOnCube(geom, randomSeed); } + else { return glm::vec3(0.0f, 0.0f, 0.0f); } +} + //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){ @@ -176,8 +245,16 @@ __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){ + + thrust::default_random_engine rng(hash(randomSeed)); + thrust::uniform_real_distribution u01(-0.5,0.5); + + float x = (float)u01(rng); + float y = (float)u01(rng); + float z = (float)u01(rng); - return glm::vec3(0,0,0); + glm::vec3 randPoint = multiplyMV(sphere.transform, glm::normalize(glm::vec4(x,y,z,1.0f))); + return randPoint; } #endif diff --git a/src/main.cpp b/src/main.cpp index 9b1fdf7..081b415 100755 --- a/src/main.cpp +++ b/src/main.cpp @@ -105,74 +105,73 @@ void runCuda(){ // Map OpenGL buffer object for writing from CUDA on a single GPU // No data is moved (Win & Linux). When mapped to CUDA, OpenGL should not use this buffer - if(iterationsiterations){ - uchar4 *dptr=NULL; - iterations++; - cudaGLMapBufferObject((void**)&dptr, pbo); + if(iterations < renderCam->iterations) { + uchar4 *dptr=NULL; + iterations++; + cudaGLMapBufferObject((void**)&dptr, pbo); - //pack geom and material arrays - geom* geoms = new geom[renderScene->objects.size()]; - material* materials = new material[renderScene->materials.size()]; + //pack geom and material arrays + geom* geoms = new geom[renderScene->objects.size()]; + material* materials = new material[renderScene->materials.size()]; - for(int i=0; iobjects.size(); i++){ - geoms[i] = renderScene->objects[i]; - } - for(int i=0; imaterials.size(); i++){ - materials[i] = renderScene->materials[i]; - } + for(int i=0; iobjects.size(); i++){ + geoms[i] = renderScene->objects[i]; + } + for(int i=0; imaterials.size(); i++){ + materials[i] = renderScene->materials[i]; + } - - // execute the kernel - cudaRaytraceCore(dptr, renderCam, targetFrame, iterations, materials, renderScene->materials.size(), geoms, renderScene->objects.size() ); + // execute the kernel + cudaRaytraceCore(dptr, renderCam, targetFrame, iterations, materials, renderScene->materials.size(), geoms, renderScene->objects.size() ); - // unmap buffer object - cudaGLUnmapBufferObject(pbo); - }else{ - - if(!finishedRender){ - //output image file - image outputImage(renderCam->resolution.x, renderCam->resolution.y); - - for(int x=0; xresolution.x; x++){ - for(int y=0; yresolution.y; y++){ - int index = x + (y * renderCam->resolution.x); - outputImage.writePixelRGB(renderCam->resolution.x-1-x,y,renderCam->image[index]); - } - } + // unmap buffer object + cudaGLUnmapBufferObject(pbo); + } + else { + + if(!finishedRender){ + //output image file + image outputImage(renderCam->resolution.x, renderCam->resolution.y); + + for(int x=0; xresolution.x; x++){ + for(int y=0; yresolution.y; y++){ + int index = x + (y * renderCam->resolution.x); + outputImage.writePixelRGB(renderCam->resolution.x-1-x,y,renderCam->image[index]); + } + } - gammaSettings gamma; - gamma.applyGamma = true; - gamma.gamma = 1.0/2.2; - gamma.divisor = renderCam->iterations; - outputImage.setGammaSettings(gamma); - string filename = renderCam->imageName; - string s; - stringstream out; - out << targetFrame; - s = out.str(); - utilityCore::replaceString(filename, ".bmp", "."+s+".bmp"); - utilityCore::replaceString(filename, ".png", "."+s+".png"); - outputImage.saveImageRGB(filename); - cout << "Saved frame " << s << " to " << filename << endl; - finishedRender = true; - if(singleFrameMode==true){ - cudaDeviceReset(); - exit(0); - } - } - if(targetFrameframes-1){ - - //clear image buffer and move onto next frame - targetFrame++; - iterations = 0; - for(int i=0; iresolution.x*renderCam->resolution.y; i++){ - renderCam->image[i] = glm::vec3(0,0,0); - } - cudaDeviceReset(); - finishedRender = false; - } - } - + gammaSettings gamma; + gamma.applyGamma = false; + gamma.gamma = 1.0/2.2; + gamma.divisor = 1.0; + outputImage.setGammaSettings(gamma); + string filename = renderCam->imageName; + string s; + stringstream out; + out << targetFrame; + s = out.str(); + utilityCore::replaceString(filename, ".bmp", "."+s+".bmp"); + utilityCore::replaceString(filename, ".png", "."+s+".png"); + outputImage.saveImageRGB(filename); + cout << "Saved frame " << s << " to " << filename << endl; + finishedRender = true; + if(singleFrameMode==true){ + cudaDeviceReset(); + exit(0); + } + } + if(targetFrame < renderCam->frames-1){ + + //clear image buffer and move onto next frame + targetFrame++; + iterations = 0; + for(int i=0; iresolution.x*renderCam->resolution.y; i++){ + renderCam->image[i] = glm::vec3(0,0,0); + } + cudaDeviceReset(); + finishedRender = false; + } + } } #ifdef __APPLE__ diff --git a/src/raytraceKernel.cu b/src/raytraceKernel.cu index b4f4ec5..160ce0a 100755 --- a/src/raytraceKernel.cu +++ b/src/raytraceKernel.cu @@ -44,9 +44,34 @@ __host__ __device__ glm::vec3 generateRandomNumberFromThread(glm::vec2 resolutio //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); + + //establish "right" camera direction + glm::normalize(eye); glm::normalize(view); + glm::vec3 right = glm::normalize(glm::cross(up, view)); + + // calculate P1 and P2 in both x and y directions + glm::vec3 image_center = eye + view; + glm::vec3 P1_X = image_center - tan((float)4.0*fov.x)*right; + glm::vec3 P2_X = image_center + tan((float)4.0*fov.x)*right; + glm::vec3 P1_Y = image_center - tan((float)4.0*fov.y)*up; + glm::vec3 P2_Y = image_center + tan((float)4.0*fov.y)*up; + + glm::vec3 bottom_left = P1_X + (P1_Y - image_center); + glm::vec3 bottom_right = P2_X + (P1_Y - image_center); + glm::vec3 top_left = P1_X + (P2_Y - image_center); + + glm::vec3 imgRight = bottom_right - bottom_left; + glm::vec3 imgUp = top_left - bottom_left; + + // supersample the pixels by taking a randomly offset ray in each iteration + glm::vec3 random_offset = generateRandomNumberFromThread(resolution, time, x, y); + float x_offset = random_offset.x; + float y_offset = random_offset.y; + glm::vec3 img_point = bottom_left + ((float)x + x_offset)/(float)resolution.x*imgRight + ((float)y + y_offset)/(float)resolution.y*imgUp; + glm::vec3 direction = glm::normalize(img_point - eye); + + // return value + ray r; r.origin = eye; r.direction = direction; return r; } @@ -97,23 +122,77 @@ __global__ void sendImageToPBO(uchar4* PBOpos, glm::vec2 resolution, glm::vec3* //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){ + staticGeom* geoms, int numberOfGeoms, staticGeom* lights, int numberOfLights, + material* materials, int iterations, int traceDepth) { - int x = (blockIdx.x * blockDim.x) + threadIdx.x; - int y = (blockIdx.y * blockDim.y) + threadIdx.y; - int index = x + (y * resolution.x); + // Find index of pixel and create empty color vector + int x = (blockIdx.x * blockDim.x) + threadIdx.x; + int y = (blockIdx.y * blockDim.y) + threadIdx.y; + int index = x + (y * resolution.x); + glm::vec3 newColor(0,0,0); + + // Get initial ray from camera through this position + ray currentRay = raycastFromCameraKernel(resolution, time, x, y, cam.position, cam.view, cam.up, cam.fov); + + ray reflectionRay; + int currentDepth = 0; + bool reflect = false; + glm::vec3 currentSpecCoeff(1.0f, 1.0f, 1.0f); + + // Return values for the intersection test + glm::vec3 intersection_point; + glm::vec3 intersection_normal; + material intersection_mtl; - if((x<=resolution.x && y<=resolution.y)){ + do { + // Find the closest geometry intersection along the ray + float t; + float min_t = -1.0; + for (int i = 0; i < numberOfGeoms; i++) { + staticGeom geom = geoms[i]; + t = geomIntersectionTest(geom, currentRay, intersection_point, intersection_normal); + if ((t > 0.0) && (t < min_t || min_t < 0.0)) { + min_t = t; + intersection_mtl = materials[geom.materialid]; + } + } + + // find reflected ray if one exists + if (intersection_mtl.hasReflective) { + reflect = true; + glm::vec3 rd = calculateReflectionDirection(intersection_normal, currentRay.direction); + glm::vec3 ro = glm::vec3(intersection_point); + reflectionRay.direction = rd; reflectionRay.origin = ro; + } + else { reflect = false; } + + // Find and clamp diffuse contribution at point + glm::vec3 phong = computePhongTotal(currentRay, intersection_point, intersection_normal, intersection_mtl, + lights, numberOfLights, geoms, numberOfGeoms, materials, (float)time); + if (phong.x > 1.0f) { phong.x = 1.0f; } else if (phong.x < 0.0f) { phong.x = 0.0f; } + if (phong.y > 1.0f) { phong.y = 1.0f; } else if (phong.y < 0.0f) { phong.y = 0.0f; } + if (phong.z > 1.0f) { phong.z = 1.0f; } else if (phong.z < 0.0f) { phong.z = 0.0f; } + newColor += (currentSpecCoeff * phong); - colors[index] = generateRandomNumberFromThread(resolution, time, x, y); - } + currentDepth++; + currentRay.origin = reflectionRay.origin; + currentRay.direction = reflectionRay.direction; + currentSpecCoeff *= intersection_mtl.specularColor; + } + while (reflect && (currentDepth < traceDepth)); + + if (newColor.x > 1.0f) { newColor.x = 1.0f; } else if (newColor.x < 0.0f) { newColor.x = 0.0f; } + if (newColor.y > 1.0f) { newColor.y = 1.0f; } else if (newColor.y < 0.0f) { newColor.y = 0.0f; } + if (newColor.z > 1.0f) { newColor.z = 1.0f; } else if (newColor.z < 0.0f) { newColor.z = 0.0f; } + if((x<=resolution.x && y<=resolution.y)) + colors[index] += newColor / (float)iterations; } //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 + int traceDepth = 3; //determines how many bounces the raytracer traces // set up crucial magic int tileSize = 8; @@ -125,10 +204,12 @@ void cudaRaytraceCore(uchar4* PBOpos, camera* renderCam, int frame, int iteratio 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); - //package geometry and materials and sent to GPU + //package geometry and send to GPU + int numberOfLights = 0; staticGeom* geomList = new staticGeom[numberOfGeoms]; for(int i=0; i 0.0) + numberOfLights++; } - + staticGeom* cudageoms = NULL; cudaMalloc((void**)&cudageoms, numberOfGeoms*sizeof(staticGeom)); cudaMemcpy( cudageoms, geomList, numberOfGeoms*sizeof(staticGeom), cudaMemcpyHostToDevice); + //package materials and send to GPU + material* materialList = new material[numberOfMaterials]; + for (int i=0; i 0.0) { + staticGeom newLight; + newLight.objectid = geoms[i].objectid; + newLight.type = geoms[i].type; + newLight.materialid = geoms[i].materialid; + newLight.translation = geoms[i].translations[frame]; + newLight.rotation = geoms[i].rotations[frame]; + newLight.scale = geoms[i].scales[frame]; + newLight.transform = geoms[i].transforms[frame]; + newLight.inverseTransform = geoms[i].inverseTransforms[frame]; + lightList[light_idx++] = newLight; + } + } + + staticGeom* cudalights = NULL; + cudaMalloc((void**)&cudalights, numberOfLights*sizeof(staticGeom)); + cudaMemcpy(cudalights, lightList, numberOfLights*sizeof(staticGeom), cudaMemcpyHostToDevice); + //package camera cameraData cam; cam.resolution = renderCam->resolution; @@ -150,19 +277,23 @@ void cudaRaytraceCore(uchar4* PBOpos, camera* renderCam, int frame, int iteratio cam.view = renderCam->views[frame]; cam.up = renderCam->ups[frame]; cam.fov = renderCam->fov; - + //kernel launches - raytraceRay<<>>(renderCam->resolution, (float)iterations, cam, traceDepth, cudaimage, cudageoms, numberOfGeoms); - + raytraceRay<<>>(renderCam->resolution, (float)iterations, cam, traceDepth, cudaimage, cudageoms, numberOfGeoms, cudalights, numberOfLights, cudamaterials, renderCam->iterations, traceDepth); + sendImageToPBO<<>>(PBOpos, renderCam->resolution, cudaimage); - + //retrieve image from GPU cudaMemcpy( renderCam->image, cudaimage, (int)renderCam->resolution.x*(int)renderCam->resolution.y*sizeof(glm::vec3), cudaMemcpyDeviceToHost); - + //free up stuff, or else we'll leak memory like a madman cudaFree( cudaimage ); cudaFree( cudageoms ); + cudaFree( cudalights ); + cudaFree( cudamaterials ); delete geomList; + delete lightList; + delete materialList; // make certain the kernel has completed cudaThreadSynchronize(); diff --git a/src/scene.cpp b/src/scene.cpp index 415d627..b194012 100755 --- a/src/scene.cpp +++ b/src/scene.cpp @@ -42,6 +42,7 @@ int scene::loadObject(string objectid){ }else{ cout << "Loading Object " << id << "..." << endl; geom newObject; + newObject.objectid = id; string line; //load object type @@ -135,8 +136,8 @@ int scene::loadObject(string objectid){ } int scene::loadCamera(){ - cout << "Loading Camera ..." << endl; - camera newCamera; + printf("Loading Camera ...\n"); + camera newCamera; float fovy; //load static properties @@ -162,6 +163,7 @@ int scene::loadCamera(){ vector positions; vector views; vector ups; + while (!line.empty() && fp_in.good()){ //check frame number diff --git a/src/sceneStructs.h b/src/sceneStructs.h index b10f1cf..790961a 100755 --- a/src/sceneStructs.h +++ b/src/sceneStructs.h @@ -20,6 +20,7 @@ struct ray { struct geom { enum GEOMTYPE type; + int objectid; int materialid; int frames; glm::vec3* translations; @@ -31,6 +32,7 @@ struct geom { struct staticGeom { enum GEOMTYPE type; + int objectid; int materialid; glm::vec3 translation; glm::vec3 rotation; diff --git a/src/utilities.cpp b/src/utilities.cpp index 3fd4b73..7ea2ea5 100755 --- a/src/utilities.cpp +++ b/src/utilities.cpp @@ -52,6 +52,25 @@ glm::vec3 utilityCore::clampRGB(glm::vec3 color){ return color; } +glm::vec3 utilityCore::clampLight(glm::vec3 color){ + if(color.x > 1.0f){ + color.x = 1.0f; + }else if(color.x < 0.0f){ + color.x = 0.0f; + } + if(color.y > 1.0f){ + color.y = 1.0f; + }else if(color.y < 0.0f){ + color.y = 0.0f; + } + if(color.z > 1.0f){ + color.z = 1.0f; + }else if(color.z < 0.0f){ + color.z = 0.0f; + } + return color; +} + bool utilityCore::epsilonCheck(float a, float b){ if(fabs(fabs(a)-fabs(b)) tokenizeString(std::string str); extern cudaMat4 glmMat4ToCudaMat4(glm::mat4 a);