diff --git a/CMakeLists.txt b/CMakeLists.txt old mode 100644 new mode 100755 index 57032a4..efa8560 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -97,4 +97,4 @@ set(SOURCE_FILES ${SOURCE_FILES} "${EXTERNAL}/src/stb_image/stb_image_write.c") cuda_add_executable(Project3-Pathtracer ${SOURCE_FILES} OPTIONS -arch=sm_20) -target_link_libraries(Project3-Pathtracer ${CORELIBS}) \ No newline at end of file +target_link_libraries(Project3-Pathtracer ${CORELIBS}) diff --git a/README.md b/README.md old mode 100644 new mode 100755 index 3b51372..6ddee5c --- a/README.md +++ b/README.md @@ -1,273 +1,55 @@ CIS 565 Project3 : CUDA Pathtracer =================== -Fall 2014 - -Due Wed, 10/8 (submit without penalty until Sun, 10/12) - -## 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 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 -* 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, called TAKUAscene. -TAKUAscene 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 TAKUAscene 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. +![](images/scene2_iterations2000.jpg) + +PROJECT DESCRIPTION +------------------- +This is a GPU CUDA version path tracer, which generate images very quickly compares to CPU version path tracer. It is based on the ray tracing algorithm by tracing rays of light through pixels in an image plane to the scene where light rays can bounce and intersect with geometries. + +Based on the given code base, first I have implemented the ray tracing part using recursive algorithm. +Reference: http://sirkan.iit.bme.hu/~szirmay/ray.pdf +![](images/1ray_iterations1000.jpg) +Other helpful reference: http://www.iquilezles.org/www/articles/simplepathtracing/simplepathtracing.htm +http://cs.brown.edu/courses/cs224/papers/mc_pathtracing.pdf + +Features Implemented +------------------- + - Simple Easy Raycasting from a camera to scene + - Diffuse surfaces + - Perfect specular reflective surfaces + - Cube intersection testing + - Sphere surface point sampling + - Stream compaction optimization + +Extra Features +------------------- + - Refraction, i.e. glass + ![](images/scene2_iterations20.jpg) + ![](images/scene2_iterations200.jpg) + ![](images/scene2_iterations2000.jpg) + + - Interactive Camera (Show in the video) + +Performance Analysis +------------------- + - Stream Compaction + Using stream compaction method to dynamically kill inactive rays and keep the wanted rays in the tracing program. The datatable and image shows the difference w/o stream compaction. + +Time(ms)/scene2,10 maxDepth iteration 50 iteration 100 iteration 150 iteration 200 iteration 250 + +with streamcompaction 9.63 21.12 34.27 47.09 58.97 + +without streamcompaction 8.66 23.56 42.75 64.51 90.32 + +![](datatable.jpg) + + - Images of the sample Scene +![](images/path_iterations100.jpg) +![](images/path_iterations1000.jpg) +![](images/path_iterations2000.jpg) + + +Vedio Link +------------------- +http://youtu.be/63mpQzdkslI diff --git a/data/scenes/sampleScene.txt b/data/scenes/sampleScene.txt old mode 100644 new mode 100755 diff --git a/data/scenes/sampleScene2.txt b/data/scenes/sampleScene2.txt new file mode 100755 index 0000000..9fda736 --- /dev/null +++ b/data/scenes/sampleScene2.txt @@ -0,0 +1,219 @@ +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 + +MATERIAL 9 //blue glossy +RGB .15 .48 .79 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 0 +REFRIOR 2.6 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 0 + +CAMERA +RES 800 800 +FOVY 25 +ITERATIONS 5000 +FILE test.bmp +FOCAL 8.2 +APERTURE 1.3 +frame 0 +EYE 0 4.5 12 +VIEW 0 0 -1 +UP 0 1 0 + +OBJECT 0 +cube +material 0 +frame 0 +TRANS 0 0 0 +ROTAT 0 0 90 +SCALE .01 10 10 + +OBJECT 1 +cube +material 0 +frame 0 +TRANS 0 5 -5 +ROTAT 0 90 0 +SCALE .01 10 10 + +OBJECT 2 +cube +material 0 +frame 0 +TRANS 0 10 0 +ROTAT 0 0 90 +SCALE .01 10 10 + +OBJECT 3 +cube +material 1 +frame 0 +TRANS -5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +OBJECT 4 +cube +material 2 +frame 0 +TRANS 5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +OBJECT 5 +cube +material 5 +frame 0 +TRANS -1 2 -1 +ROTAT 0 180 0 +SCALE 2.5 2.5 2.5 + +OBJECT 6 +sphere +material 3 +frame 0 +TRANS 2 5.4 -1 +ROTAT 0 180 0 +SCALE 2.5 2.5 2.5 + +OBJECT 7 +sphere +material 9 +frame 0 +TRANS -3.2 4.2 0 +ROTAT 0 180 0 +SCALE 2 2 2 + +OBJECT 8 +sphere +material 5 +frame 0 +TRANS 3.5 2 -1 +ROTAT 0 180 0 +SCALE 2 2 2 + +OBJECT 9 +cube +material 8 +frame 0 +TRANS 2 10 -0.5 +ROTAT 0 0 90 +SCALE .3 1 5 + +OBJECT 10 +cube +material 8 +frame 0 +TRANS -2 10 -0.5 +ROTAT 0 0 90 +SCALE .3 1 5 diff --git a/datatable.jpg b/datatable.jpg new file mode 100644 index 0000000..e766aac Binary files /dev/null and b/datatable.jpg differ diff --git a/datatable.txt b/datatable.txt new file mode 100755 index 0000000..fb0dd44 --- /dev/null +++ b/datatable.txt @@ -0,0 +1,5 @@ +scene2: (under 10 maxDepth) iteration 50 iteration 100 iteration 150 iteration 200 iteration 250 + +with streamcompaction 9.63 21.12 34.27 47.09 58.97 + +without streamcompaction 8.66 23.56 42.75 64.51 90.32 \ No newline at end of file diff --git a/external/bin/freeglut.dll b/external/bin/freeglut.dll old mode 100644 new mode 100755 diff --git a/external/bin/glew32.dll b/external/bin/glew32.dll old mode 100644 new mode 100755 diff --git a/external/bin/glut32.dll b/external/bin/glut32.dll old mode 100644 new mode 100755 diff --git a/external/include/GL/Copying.txt b/external/include/GL/Copying.txt old mode 100644 new mode 100755 diff --git a/external/include/GL/LICENSE.txt b/external/include/GL/LICENSE.txt old mode 100644 new mode 100755 diff --git a/external/include/GL/Readme.txt b/external/include/GL/Readme.txt old mode 100644 new mode 100755 diff --git a/external/include/GL/freeglut.h b/external/include/GL/freeglut.h old mode 100644 new mode 100755 diff --git a/external/include/GL/freeglut_ext.h b/external/include/GL/freeglut_ext.h old mode 100644 new mode 100755 diff --git a/external/include/GL/freeglut_std.h b/external/include/GL/freeglut_std.h old mode 100644 new mode 100755 diff --git a/external/include/GL/glew.h b/external/include/GL/glew.h old mode 100644 new mode 100755 diff --git a/external/include/GL/glut.h b/external/include/GL/glut.h old mode 100644 new mode 100755 diff --git a/external/include/GL/glxew.h b/external/include/GL/glxew.h old mode 100644 new mode 100755 diff --git a/external/include/GL/wglew.h b/external/include/GL/wglew.h old mode 100644 new mode 100755 diff --git a/external/include/GLFW/COPYING.txt b/external/include/GLFW/COPYING.txt old mode 100644 new mode 100755 diff --git a/external/include/GLFW/glfw3.h b/external/include/GLFW/glfw3.h old mode 100644 new mode 100755 diff --git a/external/include/GLFW/glfw3native.h b/external/include/GLFW/glfw3native.h old mode 100644 new mode 100755 diff --git a/external/include/glm/CMakeLists.txt b/external/include/glm/CMakeLists.txt old mode 100644 new mode 100755 diff --git a/external/include/glm/common.hpp b/external/include/glm/common.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/_features.hpp b/external/include/glm/detail/_features.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/_fixes.hpp b/external/include/glm/detail/_fixes.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/_literals.hpp b/external/include/glm/detail/_literals.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/_noise.hpp b/external/include/glm/detail/_noise.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/_swizzle.hpp b/external/include/glm/detail/_swizzle.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/_swizzle_func.hpp b/external/include/glm/detail/_swizzle_func.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/_vectorize.hpp b/external/include/glm/detail/_vectorize.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/dummy.cpp b/external/include/glm/detail/dummy.cpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/func_common.hpp b/external/include/glm/detail/func_common.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/func_common.inl b/external/include/glm/detail/func_common.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/func_exponential.hpp b/external/include/glm/detail/func_exponential.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/func_exponential.inl b/external/include/glm/detail/func_exponential.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/func_geometric.hpp b/external/include/glm/detail/func_geometric.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/func_geometric.inl b/external/include/glm/detail/func_geometric.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/func_integer.hpp b/external/include/glm/detail/func_integer.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/func_integer.inl b/external/include/glm/detail/func_integer.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/func_matrix.hpp b/external/include/glm/detail/func_matrix.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/func_matrix.inl b/external/include/glm/detail/func_matrix.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/func_noise.hpp b/external/include/glm/detail/func_noise.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/func_noise.inl b/external/include/glm/detail/func_noise.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/func_packing.hpp b/external/include/glm/detail/func_packing.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/func_packing.inl b/external/include/glm/detail/func_packing.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/func_trigonometric.hpp b/external/include/glm/detail/func_trigonometric.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/func_trigonometric.inl b/external/include/glm/detail/func_trigonometric.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/func_vector_relational.hpp b/external/include/glm/detail/func_vector_relational.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/func_vector_relational.inl b/external/include/glm/detail/func_vector_relational.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/glm.cpp b/external/include/glm/detail/glm.cpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/hint.hpp b/external/include/glm/detail/hint.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/intrinsic_common.hpp b/external/include/glm/detail/intrinsic_common.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/intrinsic_common.inl b/external/include/glm/detail/intrinsic_common.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/intrinsic_exponential.hpp b/external/include/glm/detail/intrinsic_exponential.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/intrinsic_exponential.inl b/external/include/glm/detail/intrinsic_exponential.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/intrinsic_geometric.hpp b/external/include/glm/detail/intrinsic_geometric.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/intrinsic_geometric.inl b/external/include/glm/detail/intrinsic_geometric.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/intrinsic_integer.hpp b/external/include/glm/detail/intrinsic_integer.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/intrinsic_integer.inl b/external/include/glm/detail/intrinsic_integer.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/intrinsic_matrix.hpp b/external/include/glm/detail/intrinsic_matrix.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/intrinsic_matrix.inl b/external/include/glm/detail/intrinsic_matrix.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/intrinsic_trigonometric.hpp b/external/include/glm/detail/intrinsic_trigonometric.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/intrinsic_trigonometric.inl b/external/include/glm/detail/intrinsic_trigonometric.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/intrinsic_vector_relational.hpp b/external/include/glm/detail/intrinsic_vector_relational.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/intrinsic_vector_relational.inl b/external/include/glm/detail/intrinsic_vector_relational.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/precision.hpp b/external/include/glm/detail/precision.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/precision.inl b/external/include/glm/detail/precision.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/setup.hpp b/external/include/glm/detail/setup.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/type_float.hpp b/external/include/glm/detail/type_float.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/type_gentype.hpp b/external/include/glm/detail/type_gentype.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/type_gentype.inl b/external/include/glm/detail/type_gentype.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/type_half.hpp b/external/include/glm/detail/type_half.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/type_half.inl b/external/include/glm/detail/type_half.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/type_int.hpp b/external/include/glm/detail/type_int.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/type_mat.hpp b/external/include/glm/detail/type_mat.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/type_mat.inl b/external/include/glm/detail/type_mat.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/type_mat2x2.hpp b/external/include/glm/detail/type_mat2x2.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/type_mat2x2.inl b/external/include/glm/detail/type_mat2x2.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/type_mat2x3.hpp b/external/include/glm/detail/type_mat2x3.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/type_mat2x3.inl b/external/include/glm/detail/type_mat2x3.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/type_mat2x4.hpp b/external/include/glm/detail/type_mat2x4.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/type_mat2x4.inl b/external/include/glm/detail/type_mat2x4.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/type_mat3x2.hpp b/external/include/glm/detail/type_mat3x2.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/type_mat3x2.inl b/external/include/glm/detail/type_mat3x2.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/type_mat3x3.hpp b/external/include/glm/detail/type_mat3x3.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/type_mat3x3.inl b/external/include/glm/detail/type_mat3x3.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/type_mat3x4.hpp b/external/include/glm/detail/type_mat3x4.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/type_mat3x4.inl b/external/include/glm/detail/type_mat3x4.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/type_mat4x2.hpp b/external/include/glm/detail/type_mat4x2.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/type_mat4x2.inl b/external/include/glm/detail/type_mat4x2.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/type_mat4x3.hpp b/external/include/glm/detail/type_mat4x3.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/type_mat4x3.inl b/external/include/glm/detail/type_mat4x3.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/type_mat4x4.hpp b/external/include/glm/detail/type_mat4x4.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/type_mat4x4.inl b/external/include/glm/detail/type_mat4x4.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/type_vec.hpp b/external/include/glm/detail/type_vec.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/type_vec.inl b/external/include/glm/detail/type_vec.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/type_vec1.hpp b/external/include/glm/detail/type_vec1.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/type_vec1.inl b/external/include/glm/detail/type_vec1.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/type_vec2.hpp b/external/include/glm/detail/type_vec2.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/type_vec2.inl b/external/include/glm/detail/type_vec2.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/type_vec3.hpp b/external/include/glm/detail/type_vec3.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/type_vec3.inl b/external/include/glm/detail/type_vec3.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/type_vec4.hpp b/external/include/glm/detail/type_vec4.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/detail/type_vec4.inl b/external/include/glm/detail/type_vec4.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/exponential.hpp b/external/include/glm/exponential.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/ext.hpp b/external/include/glm/ext.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/fwd.hpp b/external/include/glm/fwd.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/geometric.hpp b/external/include/glm/geometric.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/glm.hpp b/external/include/glm/glm.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtc/constants.hpp b/external/include/glm/gtc/constants.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtc/constants.inl b/external/include/glm/gtc/constants.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtc/epsilon.hpp b/external/include/glm/gtc/epsilon.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtc/epsilon.inl b/external/include/glm/gtc/epsilon.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtc/matrix_access.hpp b/external/include/glm/gtc/matrix_access.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtc/matrix_access.inl b/external/include/glm/gtc/matrix_access.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtc/matrix_integer.hpp b/external/include/glm/gtc/matrix_integer.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtc/matrix_inverse.hpp b/external/include/glm/gtc/matrix_inverse.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtc/matrix_inverse.inl b/external/include/glm/gtc/matrix_inverse.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtc/matrix_transform.hpp b/external/include/glm/gtc/matrix_transform.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtc/matrix_transform.inl b/external/include/glm/gtc/matrix_transform.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtc/noise.hpp b/external/include/glm/gtc/noise.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtc/noise.inl b/external/include/glm/gtc/noise.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtc/packing.hpp b/external/include/glm/gtc/packing.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtc/packing.inl b/external/include/glm/gtc/packing.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtc/quaternion.hpp b/external/include/glm/gtc/quaternion.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtc/quaternion.inl b/external/include/glm/gtc/quaternion.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtc/random.hpp b/external/include/glm/gtc/random.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtc/random.inl b/external/include/glm/gtc/random.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtc/reciprocal.hpp b/external/include/glm/gtc/reciprocal.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtc/reciprocal.inl b/external/include/glm/gtc/reciprocal.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtc/type_precision.hpp b/external/include/glm/gtc/type_precision.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtc/type_precision.inl b/external/include/glm/gtc/type_precision.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtc/type_ptr.hpp b/external/include/glm/gtc/type_ptr.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtc/type_ptr.inl b/external/include/glm/gtc/type_ptr.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtc/ulp.hpp b/external/include/glm/gtc/ulp.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtc/ulp.inl b/external/include/glm/gtc/ulp.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/associated_min_max.hpp b/external/include/glm/gtx/associated_min_max.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/associated_min_max.inl b/external/include/glm/gtx/associated_min_max.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/bit.hpp b/external/include/glm/gtx/bit.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/bit.inl b/external/include/glm/gtx/bit.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/closest_point.hpp b/external/include/glm/gtx/closest_point.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/closest_point.inl b/external/include/glm/gtx/closest_point.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/color_space.hpp b/external/include/glm/gtx/color_space.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/color_space.inl b/external/include/glm/gtx/color_space.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/color_space_YCoCg.hpp b/external/include/glm/gtx/color_space_YCoCg.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/color_space_YCoCg.inl b/external/include/glm/gtx/color_space_YCoCg.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/compatibility.hpp b/external/include/glm/gtx/compatibility.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/compatibility.inl b/external/include/glm/gtx/compatibility.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/component_wise.hpp b/external/include/glm/gtx/component_wise.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/component_wise.inl b/external/include/glm/gtx/component_wise.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/constants.hpp b/external/include/glm/gtx/constants.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/dual_quaternion.hpp b/external/include/glm/gtx/dual_quaternion.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/dual_quaternion.inl b/external/include/glm/gtx/dual_quaternion.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/epsilon.hpp b/external/include/glm/gtx/epsilon.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/euler_angles.hpp b/external/include/glm/gtx/euler_angles.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/euler_angles.inl b/external/include/glm/gtx/euler_angles.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/extend.hpp b/external/include/glm/gtx/extend.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/extend.inl b/external/include/glm/gtx/extend.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/extented_min_max.hpp b/external/include/glm/gtx/extented_min_max.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/extented_min_max.inl b/external/include/glm/gtx/extented_min_max.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/fast_exponential.hpp b/external/include/glm/gtx/fast_exponential.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/fast_exponential.inl b/external/include/glm/gtx/fast_exponential.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/fast_square_root.hpp b/external/include/glm/gtx/fast_square_root.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/fast_square_root.inl b/external/include/glm/gtx/fast_square_root.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/fast_trigonometry.hpp b/external/include/glm/gtx/fast_trigonometry.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/fast_trigonometry.inl b/external/include/glm/gtx/fast_trigonometry.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/gradient_paint.hpp b/external/include/glm/gtx/gradient_paint.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/gradient_paint.inl b/external/include/glm/gtx/gradient_paint.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/handed_coordinate_space.hpp b/external/include/glm/gtx/handed_coordinate_space.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/handed_coordinate_space.inl b/external/include/glm/gtx/handed_coordinate_space.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/inertia.hpp b/external/include/glm/gtx/inertia.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/inertia.inl b/external/include/glm/gtx/inertia.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/int_10_10_10_2.hpp b/external/include/glm/gtx/int_10_10_10_2.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/int_10_10_10_2.inl b/external/include/glm/gtx/int_10_10_10_2.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/integer.hpp b/external/include/glm/gtx/integer.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/integer.inl b/external/include/glm/gtx/integer.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/intersect.hpp b/external/include/glm/gtx/intersect.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/intersect.inl b/external/include/glm/gtx/intersect.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/io.hpp b/external/include/glm/gtx/io.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/io.inl b/external/include/glm/gtx/io.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/log_base.hpp b/external/include/glm/gtx/log_base.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/log_base.inl b/external/include/glm/gtx/log_base.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/matrix_cross_product.hpp b/external/include/glm/gtx/matrix_cross_product.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/matrix_cross_product.inl b/external/include/glm/gtx/matrix_cross_product.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/matrix_interpolation.hpp b/external/include/glm/gtx/matrix_interpolation.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/matrix_interpolation.inl b/external/include/glm/gtx/matrix_interpolation.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/matrix_major_storage.hpp b/external/include/glm/gtx/matrix_major_storage.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/matrix_major_storage.inl b/external/include/glm/gtx/matrix_major_storage.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/matrix_operation.hpp b/external/include/glm/gtx/matrix_operation.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/matrix_operation.inl b/external/include/glm/gtx/matrix_operation.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/matrix_query.hpp b/external/include/glm/gtx/matrix_query.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/matrix_query.inl b/external/include/glm/gtx/matrix_query.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/matrix_transform_2d.hpp b/external/include/glm/gtx/matrix_transform_2d.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/matrix_transform_2d.inl b/external/include/glm/gtx/matrix_transform_2d.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/mixed_product.hpp b/external/include/glm/gtx/mixed_product.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/mixed_product.inl b/external/include/glm/gtx/mixed_product.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/multiple.hpp b/external/include/glm/gtx/multiple.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/multiple.inl b/external/include/glm/gtx/multiple.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/noise.hpp b/external/include/glm/gtx/noise.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/norm.hpp b/external/include/glm/gtx/norm.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/norm.inl b/external/include/glm/gtx/norm.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/normal.hpp b/external/include/glm/gtx/normal.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/normal.inl b/external/include/glm/gtx/normal.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/normalize_dot.hpp b/external/include/glm/gtx/normalize_dot.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/normalize_dot.inl b/external/include/glm/gtx/normalize_dot.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/number_precision.hpp b/external/include/glm/gtx/number_precision.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/number_precision.inl b/external/include/glm/gtx/number_precision.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/optimum_pow.hpp b/external/include/glm/gtx/optimum_pow.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/optimum_pow.inl b/external/include/glm/gtx/optimum_pow.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/orthonormalize.hpp b/external/include/glm/gtx/orthonormalize.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/orthonormalize.inl b/external/include/glm/gtx/orthonormalize.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/perpendicular.hpp b/external/include/glm/gtx/perpendicular.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/perpendicular.inl b/external/include/glm/gtx/perpendicular.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/polar_coordinates.hpp b/external/include/glm/gtx/polar_coordinates.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/polar_coordinates.inl b/external/include/glm/gtx/polar_coordinates.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/projection.hpp b/external/include/glm/gtx/projection.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/projection.inl b/external/include/glm/gtx/projection.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/quaternion.hpp b/external/include/glm/gtx/quaternion.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/quaternion.inl b/external/include/glm/gtx/quaternion.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/random.hpp b/external/include/glm/gtx/random.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/raw_data.hpp b/external/include/glm/gtx/raw_data.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/raw_data.inl b/external/include/glm/gtx/raw_data.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/reciprocal.hpp b/external/include/glm/gtx/reciprocal.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/rotate_normalized_axis.hpp b/external/include/glm/gtx/rotate_normalized_axis.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/rotate_normalized_axis.inl b/external/include/glm/gtx/rotate_normalized_axis.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/rotate_vector.hpp b/external/include/glm/gtx/rotate_vector.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/rotate_vector.inl b/external/include/glm/gtx/rotate_vector.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/scalar_relational.hpp b/external/include/glm/gtx/scalar_relational.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/scalar_relational.inl b/external/include/glm/gtx/scalar_relational.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/simd_mat4.hpp b/external/include/glm/gtx/simd_mat4.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/simd_mat4.inl b/external/include/glm/gtx/simd_mat4.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/simd_quat.hpp b/external/include/glm/gtx/simd_quat.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/simd_quat.inl b/external/include/glm/gtx/simd_quat.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/simd_vec4.hpp b/external/include/glm/gtx/simd_vec4.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/simd_vec4.inl b/external/include/glm/gtx/simd_vec4.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/spline.hpp b/external/include/glm/gtx/spline.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/spline.inl b/external/include/glm/gtx/spline.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/std_based_type.hpp b/external/include/glm/gtx/std_based_type.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/std_based_type.inl b/external/include/glm/gtx/std_based_type.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/string_cast.hpp b/external/include/glm/gtx/string_cast.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/string_cast.inl b/external/include/glm/gtx/string_cast.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/transform.hpp b/external/include/glm/gtx/transform.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/transform.inl b/external/include/glm/gtx/transform.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/transform2.hpp b/external/include/glm/gtx/transform2.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/transform2.inl b/external/include/glm/gtx/transform2.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/ulp.hpp b/external/include/glm/gtx/ulp.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/unsigned_int.hpp b/external/include/glm/gtx/unsigned_int.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/unsigned_int.inl b/external/include/glm/gtx/unsigned_int.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/vec1.hpp b/external/include/glm/gtx/vec1.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/vec1.inl b/external/include/glm/gtx/vec1.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/vector_angle.hpp b/external/include/glm/gtx/vector_angle.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/vector_angle.inl b/external/include/glm/gtx/vector_angle.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/vector_query.hpp b/external/include/glm/gtx/vector_query.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/vector_query.inl b/external/include/glm/gtx/vector_query.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/wrap.hpp b/external/include/glm/gtx/wrap.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/gtx/wrap.inl b/external/include/glm/gtx/wrap.inl old mode 100644 new mode 100755 diff --git a/external/include/glm/integer.hpp b/external/include/glm/integer.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/mat2x2.hpp b/external/include/glm/mat2x2.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/mat2x3.hpp b/external/include/glm/mat2x3.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/mat2x4.hpp b/external/include/glm/mat2x4.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/mat3x2.hpp b/external/include/glm/mat3x2.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/mat3x3.hpp b/external/include/glm/mat3x3.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/mat3x4.hpp b/external/include/glm/mat3x4.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/mat4x2.hpp b/external/include/glm/mat4x2.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/mat4x3.hpp b/external/include/glm/mat4x3.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/mat4x4.hpp b/external/include/glm/mat4x4.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/matrix.hpp b/external/include/glm/matrix.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/packing.hpp b/external/include/glm/packing.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/readme.txt b/external/include/glm/readme.txt old mode 100644 new mode 100755 diff --git a/external/include/glm/trigonometric.hpp b/external/include/glm/trigonometric.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/vec2.hpp b/external/include/glm/vec2.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/vec3.hpp b/external/include/glm/vec3.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/vec4.hpp b/external/include/glm/vec4.hpp old mode 100644 new mode 100755 diff --git a/external/include/glm/vector_relational.hpp b/external/include/glm/vector_relational.hpp old mode 100644 new mode 100755 diff --git a/external/include/glslUtil/glslUtility.hpp b/external/include/glslUtil/glslUtility.hpp old mode 100644 new mode 100755 diff --git a/external/include/stb_image/stb_image.h b/external/include/stb_image/stb_image.h old mode 100644 new mode 100755 diff --git a/external/include/stb_image/stb_image_write.h b/external/include/stb_image/stb_image_write.h old mode 100644 new mode 100755 diff --git a/external/lib/linux/libGLEW.a b/external/lib/linux/libGLEW.a old mode 100644 new mode 100755 diff --git a/external/lib/linux/libglfw3.a b/external/lib/linux/libglfw3.a old mode 100644 new mode 100755 diff --git a/external/lib/osx/libGLEW.a b/external/lib/osx/libGLEW.a old mode 100644 new mode 100755 diff --git a/external/lib/osx/libglfw3.a b/external/lib/osx/libglfw3.a old mode 100644 new mode 100755 diff --git a/external/lib/win/GL/freeglut.lib b/external/lib/win/GL/freeglut.lib old mode 100644 new mode 100755 diff --git a/external/lib/win/GL/glew32.lib b/external/lib/win/GL/glew32.lib old mode 100644 new mode 100755 diff --git a/external/lib/win/GL/glew32s.lib b/external/lib/win/GL/glew32s.lib old mode 100644 new mode 100755 diff --git a/external/lib/win/GL/glut32.lib b/external/lib/win/GL/glut32.lib old mode 100644 new mode 100755 diff --git a/external/lib/win/GLFW/glfw3.lib b/external/lib/win/GLFW/glfw3.lib old mode 100644 new mode 100755 diff --git a/external/lib/win/GLFW/glfw3dll.lib b/external/lib/win/GLFW/glfw3dll.lib old mode 100644 new mode 100755 diff --git a/external/lib/win/GLFW/msvc110/glfw3.lib b/external/lib/win/GLFW/msvc110/glfw3.lib old mode 100644 new mode 100755 diff --git a/external/lib/win/GLFW/msvc110/glfw3dll.lib b/external/lib/win/GLFW/msvc110/glfw3dll.lib old mode 100644 new mode 100755 diff --git a/external/lib/win/GLFW/msvc120/glfw3.lib b/external/lib/win/GLFW/msvc120/glfw3.lib old mode 100644 new mode 100755 diff --git a/external/lib/win/GLFW/msvc120/glfw3dll.lib b/external/lib/win/GLFW/msvc120/glfw3dll.lib old mode 100644 new mode 100755 diff --git a/external/src/glslUtil/glslUtility.cpp b/external/src/glslUtil/glslUtility.cpp old mode 100644 new mode 100755 index 1f8afed..a8c362a --- a/external/src/glslUtil/glslUtility.cpp +++ b/external/src/glslUtil/glslUtility.cpp @@ -16,26 +16,22 @@ namespace glslUtility { // embedded passthrough shaders so that default passthrough shaders don't need to be loaded static std::string passthroughVS = - " #version 330 \n" - " in vec4 Position; \n" - " in vec2 Texcoords; \n" - " \n" - " out vec2 v_Texcoords; \n" + " attribute vec4 Position; \n" + " attribute vec2 Texcoords; \n" + " varying vec2 v_Texcoords; \n" " \n" " void main(void){ \n" " v_Texcoords = Texcoords; \n" " gl_Position = Position; \n" " }"; static std::string passthroughFS = - " #version 330 \n" - " in vec2 v_Texcoords; \n" + " varying vec2 v_Texcoords; \n" " \n" " uniform sampler2D u_image; \n" " \n" - " out vec4 fragColor; \n" " void main(void){ \n" - " fragColor = texture2D(u_image, v_Texcoords); \n" - " } \n"; + " gl_FragColor = texture2D(u_image, v_Texcoords); \n" + " }"; typedef struct { GLuint vertex; @@ -206,4 +202,4 @@ namespace glslUtility { return program; } -} \ No newline at end of file +} diff --git a/external/src/stb_image/stb_image.c b/external/src/stb_image/stb_image.c old mode 100644 new mode 100755 diff --git a/external/src/stb_image/stb_image_write.c b/external/src/stb_image/stb_image_write.c old mode 100644 new mode 100755 diff --git a/images/1ray_iterations1000.jpg b/images/1ray_iterations1000.jpg new file mode 100755 index 0000000..7b581a2 Binary files /dev/null and b/images/1ray_iterations1000.jpg differ diff --git a/images/path_iterations100.jpg b/images/path_iterations100.jpg new file mode 100755 index 0000000..9193ba8 Binary files /dev/null and b/images/path_iterations100.jpg differ diff --git a/images/path_iterations1000.jpg b/images/path_iterations1000.jpg new file mode 100755 index 0000000..85b72ea Binary files /dev/null and b/images/path_iterations1000.jpg differ diff --git a/images/path_iterations2000.jpg b/images/path_iterations2000.jpg new file mode 100755 index 0000000..c03354b Binary files /dev/null and b/images/path_iterations2000.jpg differ diff --git a/images/scene2_iterations20.jpg b/images/scene2_iterations20.jpg new file mode 100755 index 0000000..bccf0d2 Binary files /dev/null and b/images/scene2_iterations20.jpg differ diff --git a/images/scene2_iterations200.jpg b/images/scene2_iterations200.jpg new file mode 100755 index 0000000..426694f Binary files /dev/null and b/images/scene2_iterations200.jpg differ diff --git a/images/scene2_iterations2000.jpg b/images/scene2_iterations2000.jpg new file mode 100755 index 0000000..abf2cfc Binary files /dev/null and b/images/scene2_iterations2000.jpg differ diff --git a/src/cudaMat4.h b/src/cudaMat4.h old mode 100644 new mode 100755 diff --git a/src/image.cpp b/src/image.cpp old mode 100644 new mode 100755 diff --git a/src/image.h b/src/image.h old mode 100644 new mode 100755 diff --git a/src/interactions.h b/src/interactions.h old mode 100644 new mode 100755 index 7bf6fab..537226b --- a/src/interactions.h +++ b/src/interactions.h @@ -1,106 +1,165 @@ -// CIS565 CUDA Raytracer: A parallel raytracer for Patrick Cozzi's CIS565: GPU Computing at the University of Pennsylvania -// Written by Yining Karl Li, Copyright (c) 2012 University of Pennsylvania -// This file includes code from: -// Yining Karl Li's TAKUA Render, a massively parallel pathtracing renderer: http://www.yiningkarlli.com - -#ifndef INTERACTIONS_H -#define INTERACTIONS_H - -#include "intersections.h" - -struct Fresnel { - float reflectionCoefficient; - float transmissionCoefficient; -}; - -struct AbsorptionAndScatteringProperties{ - glm::vec3 absorptionCoefficient; - 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); -__host__ __device__ glm::vec3 calculateTransmission(glm::vec3 absorptionCoefficient, float distance); -__host__ __device__ glm::vec3 calculateTransmissionDirection(glm::vec3 normal, glm::vec3 incident, float incidentIOR, float transmittedIOR); -__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); - -// TODO (OPTIONAL): IMPLEMENT THIS FUNCTION -__host__ __device__ glm::vec3 calculateTransmission(glm::vec3 absorptionCoefficient, float distance) { - return glm::vec3(0,0,0); -} - -// 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; -} - -// 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); -} - -// 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); -} - -// 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; -} - -// 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; - - // Find a direction that is not the normal based off of whether or not the normal's components are all equal to sqrt(1/3) or whether or not at least one component is less than sqrt(1/3). Learned this trick from Peter Kutz. - - glm::vec3 directionNotNormal; - if (abs(normal.x) < SQRT_OF_ONE_THIRD) { - directionNotNormal = glm::vec3(1, 0, 0); - } else if (abs(normal.y) < SQRT_OF_ONE_THIRD) { - directionNotNormal = glm::vec3(0, 1, 0); - } else { - directionNotNormal = glm::vec3(0, 0, 1); - } - - // Use not-normal direction to generate two perpendicular directions - glm::vec3 perpendicularDirection1 = glm::normalize(glm::cross(normal, directionNotNormal)); - glm::vec3 perpendicularDirection2 = glm::normalize(glm::cross(normal, perpendicularDirection1)); - - return ( up * normal ) + ( cos(around) * over * perpendicularDirection1 ) + ( sin(around) * over * perpendicularDirection2 ); - -} - -// 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); -} - -// 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; -}; - -#endif +// CIS565 CUDA Raytracer: A parallel raytracer for Patrick Cozzi's CIS565: GPU Computing at the University of Pennsylvania +// Written by Yining Karl Li, Copyright (c) 2012 University of Pennsylvania +// This file includes code from: +// Yining Karl Li's TAKUA Render, a massively parallel pathtracing renderer: http://www.yiningkarlli.com + +#ifndef INTERACTIONS_H +#define INTERACTIONS_H + +#include "intersections.h" + +struct Fresnel { + float reflectionCoefficient; + float transmissionCoefficient; +}; + +struct AbsorptionAndScatteringProperties{ + glm::vec3 absorptionCoefficient; + float reducedScatteringCoefficient; + //constructor + AbsorptionAndScatteringProperties(glm::vec3 absorbCoeff, float scatterCoeff){ + absorptionCoefficient = absorbCoeff; + reducedScatteringCoefficient = scatterCoeff; + } +}; + +// 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); +__host__ __device__ glm::vec3 calculateTransmission(glm::vec3 absorptionCoefficient, float distance); +__host__ __device__ glm::vec3 calculateTransmissionDirection(glm::vec3 normal, glm::vec3 incident, float incidentIOR, float transmittedIOR); +__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); + +// TODO (OPTIONAL): IMPLEMENT THIS FUNCTION +__host__ __device__ glm::vec3 calculateTransmission(glm::vec3 absorptionCoefficient, float distance) { + return glm::vec3(0,0,0); +} + +// 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; +} + +// TODO (OPTIONAL): IMPLEMENT THIS FUNCTION +__host__ __device__ glm::vec3 calculateTransmissionDirection(glm::vec3 normal, glm::vec3 incident, float incidentIOR, float transmittedIOR) { + float eta = incidentIOR/transmittedIOR; + float cosI = -glm::dot(incident, normal); + float sin2T = eta * eta * (1.0f - cosI * cosI); + if(sin2T>1) + return calculateReflectionDirection(incident, normal); + return eta * incident - (eta * cosI + sqrt(1-sin2T)) * normal; +} + +// TODO (OPTIONAL): IMPLEMENT THIS FUNCTION +__host__ __device__ glm::vec3 calculateReflectionDirection(glm::vec3 normal, glm::vec3 incident) { + + glm::vec3 reflect_ray(0.0f); + if(glm::length(incident + normal)0){ //relection + return 1; + } + else if(m.hasRefractive){ //refraction + return 2; + } + else{ //diffuse + thrust::default_random_engine rng(seed); + thrust::uniform_real_distribution u01(0,1); + r.direction = calculateRandomDirectionInHemisphere( normal, (float) u01(rng), (float) u01(rng)); + r.origin = intersect + r.direction * 0.000001f; + color = m.color; + return 0; + } +}; + +#endif diff --git a/src/intersections.h b/src/intersections.h old mode 100644 new mode 100755 index c9eafb6..7eb7104 --- a/src/intersections.h +++ b/src/intersections.h @@ -1,186 +1,398 @@ -// CIS565 CUDA Raytracer: A parallel raytracer for Patrick Cozzi's CIS565: GPU Computing at the University of Pennsylvania -// Written by Yining Karl Li, Copyright (c) 2012 University of Pennsylvania -// This file includes code from: -// Yining Karl Li's TAKUA Render, a massively parallel pathtracing renderer: http://www.yiningkarlli.com - -#ifndef INTERSECTIONS_H -#define INTERSECTIONS_H - -#include -#include - -#include "sceneStructs.h" -#include "cudaMat4.h" -#include "utilities.h" - -// Some forward declarations -__host__ __device__ glm::vec3 getPointOnRay(ray r, float t); -__host__ __device__ glm::vec3 multiplyMV(cudaMat4 m, glm::vec4 v); -__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 getRandomPointOnCube(staticGeom cube, float randomSeed); - -// Handy dandy little hashing function that provides seeds for random number generation -__host__ __device__ unsigned int hash(unsigned int a){ - a = (a+0x7ed55d16) + (a<<12); - a = (a^0xc761c23c) ^ (a>>19); - a = (a+0x165667b1) + (a<<5); - a = (a+0xd3a2646c) ^ (a<<9); - a = (a+0xfd7046c5) + (a<<3); - a = (a^0xb55a4f09) ^ (a>>16); - return a; -} - -// Quick and dirty epsilon check -__host__ __device__ bool epsilonCheck(float a, float b){ - if(fabs(fabs(a)-fabs(b)) < EPSILON){ - return true; - }else{ - return false; - } -} - -// 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 -__host__ __device__ glm::vec3 multiplyMV(cudaMat4 m, glm::vec4 v){ - glm::vec3 r(1,1,1); - r.x = (m.x.x*v.x)+(m.x.y*v.y)+(m.x.z*v.z)+(m.x.w*v.w); - r.y = (m.y.x*v.x)+(m.y.y*v.y)+(m.y.z*v.z)+(m.y.w*v.w); - r.z = (m.z.x*v.x)+(m.z.y*v.y)+(m.z.z*v.z)+(m.z.w*v.w); - return r; -} - -// 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)); -} - -// 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){ - - return -1; -} - -// 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))); - - 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 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); - - 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)); - glm::vec3 xmax = multiplyMV(geom.transform, glm::vec4(.5,0,0,1)); - glm::vec3 ymax = multiplyMV(geom.transform, glm::vec4(0,.5,0,1)); - glm::vec3 zmax = multiplyMV(geom.transform, glm::vec4(0,0,.5,1)); - float xradius = glm::distance(origin, xmax); - float yradius = glm::distance(origin, ymax); - float zradius = glm::distance(origin, zmax); - 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::uniform_real_distribution u01(0,1); - thrust::uniform_real_distribution u02(-0.5,0.5); - - // Get surface areas of sides - glm::vec3 radii = getRadiuses(cube); - float side1 = radii.x * radii.y * 4.0f; //x-y face - float side2 = radii.z * radii.y * 4.0f; //y-z face - float side3 = radii.x * radii.z* 4.0f; //x-z face - float totalarea = 2.0f * (side1+side2+side3); - - // Pick random face, weighted by surface area - float russianRoulette = (float)u01(rng); - - glm::vec3 point = glm::vec3(.5,.5,.5); - - if(russianRoulette<(side1/totalarea)){ - // x-y face - point = glm::vec3((float)u02(rng), (float)u02(rng), .5); - }else if(russianRoulette<((side1*2)/totalarea)){ - // x-y-back face - point = glm::vec3((float)u02(rng), (float)u02(rng), -.5); - }else if(russianRoulette<(((side1*2)+(side2))/totalarea)){ - // y-z face - point = glm::vec3(.5, (float)u02(rng), (float)u02(rng)); - }else if(russianRoulette<(((side1*2)+(side2*2))/totalarea)){ - // y-z-back face - point = glm::vec3(-.5, (float)u02(rng), (float)u02(rng)); - }else if(russianRoulette<(((side1*2)+(side2*2)+(side3))/totalarea)){ - // x-z face - point = glm::vec3((float)u02(rng), .5, (float)u02(rng)); - }else{ - // x-z-back face - point = glm::vec3((float)u02(rng), -.5, (float)u02(rng)); - } - - glm::vec3 randPoint = multiplyMV(cube.transform, glm::vec4(point,1.0f)); - - return randPoint; - -} - -// TODO: IMPLEMENT THIS FUNCTION -// Generates a random point on a given sphere -__host__ __device__ glm::vec3 getRandomPointOnSphere(staticGeom sphere, float randomSeed){ - - return glm::vec3(0,0,0); -} - -#endif - - +// CIS565 CUDA Raytracer: A parallel raytracer for Patrick Cozzi's CIS565: GPU Computing at the University of Pennsylvania +// Written by Yining Karl Li, Copyright (c) 2012 University of Pennsylvania +// This file includes code from: +// Yining Karl Li's TAKUA Render, a massively parallel pathtracing renderer: http://www.yiningkarlli.com + +#ifndef INTERSECTIONS_H +#define INTERSECTIONS_H + +#include +#include +#include +#include "sceneStructs.h" +#include "cudaMat4.h" +#include "utilities.h" + +// Some forward declarations +__host__ __device__ glm::vec3 getPointOnRay(ray r, float t); +__host__ __device__ glm::vec3 multiplyMV(cudaMat4 m, glm::vec4 v); +__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 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){ + a = (a+0x7ed55d16) + (a<<12); + a = (a^0xc761c23c) ^ (a>>19); + a = (a+0x165667b1) + (a<<5); + a = (a+0xd3a2646c) ^ (a<<9); + a = (a+0xfd7046c5) + (a<<3); + a = (a^0xb55a4f09) ^ (a>>16); + return a; +} + +// Quick and dirty epsilon check +__host__ __device__ bool epsilonCheck(float a, float b){ + if(fabs(fabs(a)-fabs(b)) < EPSILON){ + return true; + }else{ + return false; + } +} + +// 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 +__host__ __device__ glm::vec3 multiplyMV(cudaMat4 m, glm::vec4 v){ + glm::vec3 r(1,1,1); + r.x = (m.x.x*v.x)+(m.x.y*v.y)+(m.x.z*v.z)+(m.x.w*v.w); + r.y = (m.y.x*v.x)+(m.y.y*v.y)+(m.y.z*v.z)+(m.y.w*v.w); + r.z = (m.z.x*v.x)+(m.z.y*v.y)+(m.z.z*v.z)+(m.z.w*v.w); + return r; +} + + +// 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)); +} + +//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){ + ray rt; + + glm::vec3 bounds[2]; + if(box.type == MESH){ + bounds[0] = box.boundingBoxMin; + bounds[1] = box.boundingBoxMax; + rt.origin = r.origin; + rt.direction = r.direction; //already transformed + } + else{ + rt.origin = multiplyMV ( box.inverseTransform, glm::vec4( r.origin, 1.0f )); + rt.direction = glm::normalize( multiplyMV ( box.inverseTransform, glm::vec4( r.direction, 0.0f )) ); + bounds[0] = glm::vec3(-0.5f, -0.5f, -0.5f); + bounds[1] = glm::vec3(0.5f, 0.5f, 0.5f); + } + + glm::vec3 invDir = getInverseDirectionOfRay(rt); + int sign[3]; + sign[0] = (invDir.x < 0); + sign[1] = (invDir.y < 0); + sign[2] = (invDir.z < 0); + + double tmin, tmax, tymin, tymax, tzmin, tzmax, t; + tmin = (bounds[sign[0]].x - rt.origin.x) * invDir.x; + tmax = (bounds[1-sign[0]].x - rt.origin.x) * invDir.x; + tymin = (bounds[sign[1]].y - rt.origin.y) * invDir.y; + tymax = (bounds[1-sign[1]].y - rt.origin.y) * invDir.y; + + if ((tmin > tymax) || (tymin > tmax)) + return -1.0f; + if (tymin > tmin) + tmin = tymin; + if (tymax < tmax) + tmax = tymax; + tzmin = (bounds[(int)sign[2]].z - rt.origin.z) * invDir.z; + tzmax = (bounds[1-(int)sign[2]].z - rt.origin.z) * invDir.z; + + if ((tmin > tzmax) || (tzmin > tmax)) + return -1.0f; + if (tzmin > tmin) + tmin = tzmin; + if (tzmax < tmax) + tmax = tzmax; + if (tmax < 0.0) + return -1.0f; + if (tmin < 0.0) + t = tmax; // inside + else + t = tmin; //outside + + glm::vec3 intersect_os = float(t) * rt.direction + rt.origin; + glm::vec3 tempNormal; + if( fabs(intersect_os.x - bounds[0].x) < EPSILON) + tempNormal = glm::vec3(-1.0f, 0.0f, 0.0f); + else if( fabs( intersect_os.x - bounds[1].x) < EPSILON) + tempNormal = glm::vec3(1.0f, 0.0f, 0.0f); + else if( fabs( intersect_os.y - bounds[0].y) < EPSILON) + tempNormal = glm::vec3(0.0f, -1.0f, 0.0f); + else if( fabs( intersect_os.y - bounds[1].y) < EPSILON) + tempNormal = glm::vec3(0.0f, 1.0f, 0.0f); + else if( fabs( intersect_os.z - bounds[0].z) < EPSILON) + tempNormal = glm::vec3(0.0f, 0.0f, -1.0f); + else if( fabs( intersect_os.z - bounds[1].z) < EPSILON) + tempNormal = glm::vec3(0.0f, 0.0f, 1.0f); + + intersectionPoint = multiplyMV(box.transform, glm::vec4(intersect_os, 1.0f)); + + normal = glm::normalize(multiplyMV(box.transform, glm::vec4(tempNormal, 0.0f))); + + return glm::length(r.origin - intersectionPoint); + + //a slow traditional version of box intersection test + /*cudaMat4 Ti = box.inverseTransform; + glm::vec3 R0 = multiplyMV(Ti, glm::vec4(r.origin, 1.0)); + float l = glm::length(r.direction); + glm::vec3 Rd = multiplyMV(Ti, glm::vec4(glm::normalize(r.direction), 0.0)); + double tnear = -10000000; + double tfar = 10000000; + int slab = 0; + double t, t1, t2; + while(slab < 3){ + if(Rd[slab] == 0){ + if(R0[slab] > .5 || R0[slab] < -.5){ + return -1; + } + } + t1 = (-.5 - R0[slab]) / Rd[slab]; + t2 = (.5 - R0[slab]) / Rd[slab]; + if(t1 > t2){ + double temp = t1; + t1 = t2; + t2 = temp; + } + if(t1 > tnear) tnear = t1; + if(t2 < tfar) tfar = t2; + if(tnear > tfar){ + return -1; + } + if(tfar < 0){ + return -1; + } + slab++; + } + + if(tnear > -.0001) t = tnear; + else t = tfar; + + glm::vec3 p = R0 + (float)t * Rd; + glm::vec3 realIntersectionPoint = multiplyMV(box.transform, glm::vec4(p, 1.0)); + intersectionPoint = realIntersectionPoint; + + glm::vec4 temp_normal; + if(abs(p[0] - .5) < .001){ + temp_normal = glm::vec4(1,0,0,0); + }else if(abs(p[0] + .5) < .001){ + temp_normal = glm::vec4(-1,0,0,0); + }else if(abs(p[1] - .5) < .001){ + temp_normal = glm::vec4(0,1,0,0); + }else if(abs(p[1] + .5) < .001){ + temp_normal = glm::vec4(0,-1,0,0); + }else if(abs(p[2] - .5) < .001){ + temp_normal = glm::vec4(0,0,1,0); + }else if(abs(p[2] + .5) < .001){ + temp_normal = glm::vec4(0,0,-1,0); + } + normal = glm::normalize(multiplyMV(box.transform, temp_normal)); + + 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. +// 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))); + + 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 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); + + return glm::length(r.origin - realIntersectionPoint); +} + +__host__ __device__ float triangleIntersectionTest(triangle& tri, ray rt, glm::vec3& intersectionPoint, glm::vec3& normal){ + + glm::vec3 p1 = tri.p1; + glm::vec3 p2 = tri.p2; + glm::vec3 p3 = tri.p3; + glm::vec3 edge1 = p2 - p1; + glm::vec3 edge2 = p3 - p1; + glm::vec3 n = glm::cross(edge1, edge2); + + if (glm::length(n) < EPSILON) + return -1; + + glm::vec3 w0 = rt.origin - p1; + float a = - glm::dot(n, w0); + float b = glm::dot(n, rt.direction); + if (fabs(b) < EPSILON) + return -1; + if (fabs(a) < EPSILON) + return -1; + + float L = a/b; + if (L < 0) + return -1; + + glm::vec3 p = rt.origin + L*rt.direction; + + float e11 = glm::dot(edge1, edge1); + float e12 = glm::dot(edge1, edge2); + float e22 = glm::dot(edge2, edge2); + glm::vec3 w = p - p1; + float we1 = glm::dot(w, edge1); + float we2 = glm::dot(w, edge2); + + float D = e12 * e12 - e11 * e22; + float s = (e12 * we2 - e22 * we1) / D; + float t = (e12 * we1 - e11 * we2) / D; + + if(s < 0 || s > 1 || t < 0 || (s + t) > 1) + return -1; + + normal = glm::normalize(n); + intersectionPoint = p; + return L; + +} + +__host__ __device__ float polygonIntersectionTest(staticGeom polygon, ray r, glm::vec3& intersectionPoint, glm::vec3& normal, triangle * cudatris){ + + glm::vec3 P = multiplyMV (polygon.inverseTransform, glm::vec4(r.origin, 1.0f)); + glm::vec3 V = glm::normalize( multiplyMV (polygon.inverseTransform, glm::vec4(r.direction, 0.0f))); + ray rt; rt.origin = P; rt.direction = V; + float t = -1, tMin = -2; + glm::vec3 localIntersect, localNormal; + if(boxIntersectionTest(polygon, rt, intersectionPoint, normal) > 0){ + glm::vec3 tempIntersect, tempNormal; + + for (int i=0; i -0.5f ) || ( tMin > -1 && t < tMin && t > -0.5f ) ){ + tMin = t; + localIntersect = tempIntersect; + localNormal = tempNormal; + } + } + intersectionPoint = multiplyMV(polygon.transform, glm::vec4(localIntersect, 1.0f)); + glm::vec3 normalP_OS = intersectionPoint + localNormal; + glm::vec3 normalP_WS = multiplyMV(polygon.transform, glm::vec4(normalP_OS, 1.0f)); + normal = glm::normalize(normalP_WS - intersectionPoint); + return tMin; + } + return -1; +} + +// 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)); + glm::vec3 xmax = multiplyMV(geom.transform, glm::vec4(.5,0,0,1)); + glm::vec3 ymax = multiplyMV(geom.transform, glm::vec4(0,.5,0,1)); + glm::vec3 zmax = multiplyMV(geom.transform, glm::vec4(0,0,.5,1)); + float xradius = glm::distance(origin, xmax); + float yradius = glm::distance(origin, ymax); + float zradius = glm::distance(origin, zmax); + 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::uniform_real_distribution u01(0,1); + thrust::uniform_real_distribution u02(-0.5,0.5); + + // Get surface areas of sides + glm::vec3 radii = getRadiuses(cube); + float side1 = radii.x * radii.y * 4.0f; //x-y face + float side2 = radii.z * radii.y * 4.0f; //y-z face + float side3 = radii.x * radii.z* 4.0f; //x-z face + float totalarea = 2.0f * (side1+side2+side3); + + // Pick random face, weighted by surface area + float russianRoulette = (float)u01(rng); + + glm::vec3 point = glm::vec3(.5,.5,.5); + + if(russianRoulette<(side1/totalarea)){ + // x-y face + point = glm::vec3((float)u02(rng), (float)u02(rng), .5); + }else if(russianRoulette<((side1*2)/totalarea)){ + // x-y-back face + point = glm::vec3((float)u02(rng), (float)u02(rng), -.5); + }else if(russianRoulette<(((side1*2)+(side2))/totalarea)){ + // y-z face + point = glm::vec3(.5, (float)u02(rng), (float)u02(rng)); + }else if(russianRoulette<(((side1*2)+(side2*2))/totalarea)){ + // y-z-back face + point = glm::vec3(-.5, (float)u02(rng), (float)u02(rng)); + }else if(russianRoulette<(((side1*2)+(side2*2)+(side3))/totalarea)){ + // x-z face + point = glm::vec3((float)u02(rng), .5, (float)u02(rng)); + }else{ + // x-z-back face + point = glm::vec3((float)u02(rng), -.5, (float)u02(rng)); + } + + glm::vec3 randPoint = multiplyMV(cube.transform, glm::vec4(point,1.0f)); + + return randPoint; + +} + +// 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(-PI,PI); + glm::vec3 radius = getRadiuses(sphere); + float theta = (float)u01(rng); + float phi = (float)u01(rng); + glm::vec3 p = glm::vec3(glm::sin(theta) * glm::cos(phi), glm::sin(theta) * glm::sin(phi), glm::cos(theta)); + glm::vec3 randPoint = multiplyMV(sphere.transform, glm::vec4(p, 1.0)); + return randPoint; + + //return glm::vec3(0,0,0); +} + + +#endif + + diff --git a/src/main.cpp b/src/main.cpp old mode 100644 new mode 100755 index 7a1098f..0228147 --- a/src/main.cpp +++ b/src/main.cpp @@ -1,332 +1,473 @@ -// CIS565 CUDA Raytracer: A parallel raytracer for Patrick Cozzi's CIS565: GPU Computing at the University of Pennsylvania -// Written by Yining Karl Li, Copyright (c) 2012 University of Pennsylvania -// This file includes code from: -// Rob Farber for CUDA-GL interop, from CUDA Supercomputing For The Masses: http://www.drdobbs.com/architecture-and-design/cuda-supercomputing-for-the-masses-part/222600097 -// Varun Sampath and Patrick Cozzi for GLSL Loading, from CIS565 Spring 2012 HW5 at the University of Pennsylvania: http://cis565-spring-2012.github.com/ -// Yining Karl Li's TAKUA Render, a massively parallel pathtracing renderer: http://www.yiningkarlli.com - -#include "main.h" -#define GLEW_STATIC - -//------------------------------- -//-------------MAIN-------------- -//------------------------------- - -int main(int argc, char** argv){ - #ifdef __APPLE__ - // Needed in OSX to force use of OpenGL3.2 - glfwWindowHint(GLFW_OPENGL_VERSION_MAJOR, 3); - glfwWindowHint(GLFW_OPENGL_VERSION_MINOR, 2); - glfwWindowHint(GLFW_OPENGL_FORWARD_COMPAT, GL_TRUE); - glfwWindowHint(GLFW_OPENGL_PROFILE, GLFW_OPENGL_CORE_PROFILE); - #endif - - // Set up pathtracer stuff - bool loadedScene = false; - finishedRender = false; - - targetFrame = 0; - singleFrameMode = false; - - // Load scene file - for(int i=1; irenderCam; - width = renderCam->resolution[0]; - height = renderCam->resolution[1]; - - if(targetFrame >= renderCam->frames){ - cout << "Warning: Specified target frame is out of range, defaulting to frame 0." << endl; - targetFrame = 0; - } - - // Initialize CUDA and GL components - if (init(argc, argv)) { - // GLFW main loop - mainLoop(); - } - - return 0; -} - -void mainLoop() { - while(!glfwWindowShouldClose(window)){ - glfwPollEvents(); - runCuda(); - - string title = "CIS565 Render | " + utilityCore::convertIntToString(iterations) + " Iterations"; - glfwSetWindowTitle(window, title.c_str()); - - glBindBuffer(GL_PIXEL_UNPACK_BUFFER, pbo); - glBindTexture(GL_TEXTURE_2D, displayImage); - glTexSubImage2D(GL_TEXTURE_2D, 0, 0, 0, width, height, GL_RGBA, GL_UNSIGNED_BYTE, NULL); - glClear(GL_COLOR_BUFFER_BIT); - - // VAO, shader program, and texture already bound - glDrawElements(GL_TRIANGLES, 6, GL_UNSIGNED_SHORT, 0); - glfwSwapBuffers(window); - } - glfwDestroyWindow(window); - glfwTerminate(); -} - -//------------------------------- -//---------RUNTIME STUFF--------- -//------------------------------- - -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(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()]; - - for (int i=0; i < renderScene->objects.size(); i++) { - geoms[i] = renderScene->objects[i]; - } - for (int i=0; i < renderScene->materials.size(); i++) { - materials[i] = renderScene->materials[i]; - } - - // 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; 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]); - } - } - - gammaSettings gamma; - gamma.applyGamma = true; - gamma.gamma = 1.0; - 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; - } - } -} - -//------------------------------- -//----------SETUP STUFF---------- -//------------------------------- - -bool init(int argc, char* argv[]) { - glfwSetErrorCallback(errorCallback); - - if (!glfwInit()) { - return false; - } - - width = 800; - height = 800; - window = glfwCreateWindow(width, height, "CIS 565 Pathtracer", NULL, NULL); - if (!window){ - glfwTerminate(); - return false; - } - glfwMakeContextCurrent(window); - glfwSetKeyCallback(window, keyCallback); - - // Set up GL context - glewExperimental = GL_TRUE; - if(glewInit()!=GLEW_OK){ - return false; - } - - // Initialize other stuff - initVAO(); - initTextures(); - initCuda(); - initPBO(); - - GLuint passthroughProgram; - passthroughProgram = initShader(); - - glUseProgram(passthroughProgram); - glActiveTexture(GL_TEXTURE0); - - return true; -} - -void initPBO(){ - // set up vertex data parameter - int num_texels = width*height; - int num_values = num_texels * 4; - int size_tex_data = sizeof(GLubyte) * num_values; - - // Generate a buffer ID called a PBO (Pixel Buffer Object) - glGenBuffers(1, &pbo); - - // Make this the current UNPACK buffer (OpenGL is state-based) - glBindBuffer(GL_PIXEL_UNPACK_BUFFER, pbo); - - // Allocate data for the buffer. 4-channel 8-bit image - glBufferData(GL_PIXEL_UNPACK_BUFFER, size_tex_data, NULL, GL_DYNAMIC_COPY); - cudaGLRegisterBufferObject(pbo); - -} - -void initCuda(){ - // Use device with highest Gflops/s - cudaGLSetGLDevice(0); - - // Clean up on program exit - atexit(cleanupCuda); -} - -void initTextures(){ - glGenTextures(1, &displayImage); - glBindTexture(GL_TEXTURE_2D, displayImage); - glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST); - glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST); - glTexImage2D( GL_TEXTURE_2D, 0, GL_RGBA8, width, height, 0, GL_BGRA, GL_UNSIGNED_BYTE, NULL); -} - -void initVAO(void){ - GLfloat vertices[] = - { - -1.0f, -1.0f, - 1.0f, -1.0f, - 1.0f, 1.0f, - -1.0f, 1.0f, - }; - - GLfloat texcoords[] = - { - 1.0f, 1.0f, - 0.0f, 1.0f, - 0.0f, 0.0f, - 1.0f, 0.0f - }; - - GLushort indices[] = { 0, 1, 3, 3, 1, 2 }; - - GLuint vertexBufferObjID[3]; - glGenBuffers(3, vertexBufferObjID); - - glBindBuffer(GL_ARRAY_BUFFER, vertexBufferObjID[0]); - glBufferData(GL_ARRAY_BUFFER, sizeof(vertices), vertices, GL_STATIC_DRAW); - glVertexAttribPointer((GLuint)positionLocation, 2, GL_FLOAT, GL_FALSE, 0, 0); - glEnableVertexAttribArray(positionLocation); - - glBindBuffer(GL_ARRAY_BUFFER, vertexBufferObjID[1]); - glBufferData(GL_ARRAY_BUFFER, sizeof(texcoords), texcoords, GL_STATIC_DRAW); - glVertexAttribPointer((GLuint)texcoordsLocation, 2, GL_FLOAT, GL_FALSE, 0, 0); - glEnableVertexAttribArray(texcoordsLocation); - - glBindBuffer(GL_ELEMENT_ARRAY_BUFFER, vertexBufferObjID[2]); - glBufferData(GL_ELEMENT_ARRAY_BUFFER, sizeof(indices), indices, GL_STATIC_DRAW); -} - -GLuint initShader() { - const char *attribLocations[] = { "Position", "Tex" }; - GLuint program = glslUtility::createDefaultProgram(attribLocations, 2); - GLint location; - - glUseProgram(program); - if ((location = glGetUniformLocation(program, "u_image")) != -1) - { - glUniform1i(location, 0); - } - - return program; -} - -//------------------------------- -//---------CLEANUP STUFF--------- -//------------------------------- - -void cleanupCuda(){ - if(pbo) deletePBO(&pbo); - if(displayImage) deleteTexture(&displayImage); -} - -void deletePBO(GLuint* pbo){ - if (pbo) { - // unregister this buffer object with CUDA - cudaGLUnregisterBufferObject(*pbo); - - glBindBuffer(GL_ARRAY_BUFFER, *pbo); - glDeleteBuffers(1, pbo); - - *pbo = (GLuint)NULL; - } -} - -void deleteTexture(GLuint* tex){ - glDeleteTextures(1, tex); - *tex = (GLuint)NULL; -} - -//------------------------------ -//-------GLFW CALLBACKS--------- -//------------------------------ - -void errorCallback(int error, const char* description){ - fputs(description, stderr); -} - -void keyCallback(GLFWwindow* window, int key, int scancode, int action, int mods){ - if(key == GLFW_KEY_ESCAPE && action == GLFW_PRESS){ - glfwSetWindowShouldClose(window, GL_TRUE); - } -} +// CIS565 CUDA Raytracer: A parallel raytracer for Patrick Cozzi's CIS565: GPU Computing at the University of Pennsylvania +// Written by Yining Karl Li, Copyright (c) 2012 University of Pennsylvania +// This file includes code from: +// Rob Farber for CUDA-GL interop, from CUDA Supercomputing For The Masses: http://www.drdobbs.com/architecture-and-design/cuda-supercomputing-for-the-masses-part/222600097 +// Varun Sampath and Patrick Cozzi for GLSL Loading, from CIS565 Spring 2012 HW5 at the University of Pennsylvania: http://cis565-spring-2012.github.com/ +// Yining Karl Li's TAKUA Render, a massively parallel pathtracing renderer: http://www.yiningkarlli.com + + +#include "main.h" +#define GLEW_STATIC +#define CAM_MOVE 0.3f + + +//------------------------------- +//-------------MAIN-------------- +//------------------------------- +/*void mouseClick(int button, int status,int x,int y); +void mouseMotion(int x, int y); +static bool r_buttonDown = false; +static bool l_buttonDown = false; +static int g_yclick = 0; +static int g_ylclick = 0; +static int g_xlclick = 0; +int mouse_old_x, mouse_old_y; +glm::vec3 originPos; +glm::vec3 originView; +*/ +float fps = 0; +float preTime = 0; +float currTime; +int frames = 0; +void CalcFPS() +{ + frames ++; + currTime = glfwGetTime(); + if(currTime - preTime > 1000) + { + fps = frames*1000.0/(currTime - preTime); + preTime = currTime; + frames = 0; + } +} + +//------------------------------- +//-------------MAIN-------------- +//------------------------------- + +int main(int argc, char** argv){ + + #ifdef __APPLE__ + // Needed in OSX to force use of OpenGL3.2 + glfwWindowHint(GLFW_OPENGL_VERSION_MAJOR, 3); + glfwWindowHint(GLFW_OPENGL_VERSION_MINOR, 2); + glfwWindowHint(GLFW_OPENGL_FORWARD_COMPAT, GL_TRUE); + glfwWindowHint(GLFW_OPENGL_PROFILE, GLFW_OPENGL_CORE_PROFILE); + #endif + + // Set up pathtracer stuff + bool loadedScene = true; + finishedRender = false; + + targetFrame = 0; + singleFrameMode = false; + + // Load scene file + for(int i=1; irenderCam; + width = renderCam->resolution[0]; + height = renderCam->resolution[1]; + + if(targetFrame >= renderCam->frames){ + cout << "Warning: Specified target frame is out of range , defaulting to frame 0." << endl; + targetFrame = 0; + } + + // Initialize CUDA and GL components + if (init(argc, argv)) { + // GLFW main loop + mainLoop(); + } + + system("PAUSE"); + return 0; +} + + +void mainLoop() { + while(!glfwWindowShouldClose(window)){ + glfwPollEvents(); + + runCuda(); + //CalcFPS(); + //printf("fps: %f", fps); + string title = "GPU Pathtracer | Iterations: " + utilityCore::convertIntToString(iterations); /*+ " | FPS: " + utilityCore::convertFloatToString(fps);*/ + glfwSetWindowTitle(window, title.c_str()); + + glBindBuffer(GL_PIXEL_UNPACK_BUFFER, pbo); + glBindTexture(GL_TEXTURE_2D, displayImage); + glTexSubImage2D(GL_TEXTURE_2D, 0, 0, 0, width, height, GL_RGBA, GL_UNSIGNED_BYTE, NULL); + glClear(GL_COLOR_BUFFER_BIT); + + // VAO, shader program, and texture already bound + glDrawElements(GL_TRIANGLES, 6, GL_UNSIGNED_SHORT, 0); + glfwSwapBuffers(window); + } + glfwDestroyWindow(window); + glfwTerminate(); +} + +//------------------------------- +//---------RUNTIME STUFF--------- +//------------------------------- + +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(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()]; + + + for (int i=0; i < renderScene->objects.size(); i++) { + geoms[i] = renderScene->objects[i]; + } + for (int i=0; i < renderScene->materials.size(); i++) { + materials[i] = renderScene->materials[i]; + } + + //cout<<"iterations#: "<materials.size(), geoms, renderScene->objects.size()); + + // execute the kernel for path tracing + cudaPathTraceCore(dptr, renderCam, targetFrame, iterations, materials, renderScene->materials.size(), geoms, renderScene->objects.size()); + + // unmap buffer object + cudaGLUnmapBufferObject(pbo); + } else { + cout<<"2"<resolution.x, renderCam->resolution.y); + + 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]); + } + } + + gammaSettings gamma; + gamma.applyGamma = true; + gamma.gamma = 1.0; + gamma.divisor = 1.0; + //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 (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; + } + } +} + +//------------------------------- +//----------SETUP STUFF---------- +//------------------------------- + +bool init(int argc, char* argv[]) { + glfwSetErrorCallback(errorCallback); + + if (!glfwInit()) { + return false; + } + + width = 800; + height = 800; + window = glfwCreateWindow(width, height, "GPU Pathtracer", NULL, NULL); + if (!window){ + glfwTerminate(); + return false; + } + glfwMakeContextCurrent(window); + glfwSetKeyCallback(window, keyCallback); + //glfwSetMouseButtonCallback(window, mouseClick); + + // Set up GL context + glewExperimental = GL_TRUE; + if(glewInit()!=GLEW_OK){ + return false; + } + + // Initialize other stuff + initVAO(); + initTextures(); + initCuda(); + initPBO(); + + + GLuint passthroughProgram; + passthroughProgram = initShader(); + + glUseProgram(passthroughProgram); + glActiveTexture(GL_TEXTURE0); + + return true; +} + +void initPBO(){ + // set up vertex data parameter + int num_texels = width*height; + int num_values = num_texels * 4; + int size_tex_data = sizeof(GLubyte) * num_values; + + // Generate a buffer ID called a PBO (Pixel Buffer Object) + glGenBuffers(1, &pbo); + + // Make this the current UNPACK buffer (OpenGL is state-based) + glBindBuffer(GL_PIXEL_UNPACK_BUFFER, pbo); + + // Allocate data for the buffer. 4-channel 8-bit image + glBufferData(GL_PIXEL_UNPACK_BUFFER, size_tex_data, NULL, GL_DYNAMIC_COPY); + cudaGLRegisterBufferObject(pbo); + +} + +void initCuda(){ + + cudaGLSetGLDevice(0); + + // Clean up on program exit + atexit(cleanupCuda); + +} + +void initTextures(){ + glGenTextures(1, &displayImage); + glBindTexture(GL_TEXTURE_2D, displayImage); + glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST); + glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST); + glTexImage2D( GL_TEXTURE_2D, 0, GL_RGBA8, width, height, 0, GL_BGRA, GL_UNSIGNED_BYTE, NULL); +} + +void initVAO(void){ + GLfloat vertices[] = + { + -1.0f, -1.0f, + 1.0f, -1.0f, + 1.0f, 1.0f, + -1.0f, 1.0f, + }; + + GLfloat texcoords[] = + { + 1.0f, 1.0f, + 0.0f, 1.0f, + 0.0f, 0.0f, + 1.0f, 0.0f + }; + + GLushort indices[] = { 0, 1, 3, 3, 1, 2 }; + + GLuint vertexBufferObjID[3]; + glGenBuffers(3, vertexBufferObjID); + + glBindBuffer(GL_ARRAY_BUFFER, vertexBufferObjID[0]); + glBufferData(GL_ARRAY_BUFFER, sizeof(vertices), vertices, GL_STATIC_DRAW); + glVertexAttribPointer((GLuint)positionLocation, 2, GL_FLOAT, GL_FALSE, 0, 0); + glEnableVertexAttribArray(positionLocation); + + glBindBuffer(GL_ARRAY_BUFFER, vertexBufferObjID[1]); + glBufferData(GL_ARRAY_BUFFER, sizeof(texcoords), texcoords, GL_STATIC_DRAW); + glVertexAttribPointer((GLuint)texcoordsLocation, 2, GL_FLOAT, GL_FALSE, 0, 0); + glEnableVertexAttribArray(texcoordsLocation); + + glBindBuffer(GL_ELEMENT_ARRAY_BUFFER, vertexBufferObjID[2]); + glBufferData(GL_ELEMENT_ARRAY_BUFFER, sizeof(indices), indices, GL_STATIC_DRAW); +} + +GLuint initShader() { + const char *attribLocations[] = { "Position", "Texcoords" }; + GLuint program = glslUtility::createDefaultProgram(attribLocations, 2); + GLint location; + + glUseProgram(program); + if ((location = glGetUniformLocation(program, "u_image")) != -1) + { + glUniform1i(location, 0); + } + + return program; +} + + +void initMesh(){ //intialize cuda memory for the triangle list in MESH objects + + for(int i = 0; i < renderScene->objects.size(); i++){ + if(renderScene->objects[i].type == MESH){ + triangle * cudatris = NULL; + cudaMalloc((void**)&cudatris, renderScene->objects[i].numOfTris *sizeof(triangle)); + cudaMemcpy( cudatris, renderScene->objects[i].tris, renderScene->objects[i].numOfTris *sizeof(triangle), cudaMemcpyHostToDevice); + renderScene->objects[i].tris = cudatris; + } + } + + +} + +//------------------------------- +//---------CLEANUP STUFF--------- +//------------------------------- + +void cleanupCuda(){ + if(pbo) deletePBO(&pbo); + if(displayImage) deleteTexture(&displayImage); +} + +void deletePBO(GLuint* pbo){ + if (pbo) { + // unregister this buffer object with CUDA + cudaGLUnregisterBufferObject(*pbo); + + glBindBuffer(GL_ARRAY_BUFFER, *pbo); + glDeleteBuffers(1, pbo); + + *pbo = (GLuint)NULL; + } +} + +void deleteTexture(GLuint* tex){ + glDeleteTextures(1, tex); + *tex = (GLuint)NULL; +} + +//------------------------------ +//-------GLFW CALLBACKS--------- +//------------------------------ + +void errorCallback(int error, const char* description){ + fputs(description, stderr); +} + +void cameraReset(){ + iterations = 0; + for(int i = 0; i < width * height; i++) + renderCam->image[i] = glm::vec3(0.0f); +} + +void keyCallback(GLFWwindow* window, int key, int scancode, int action, int mods){ + if(key == GLFW_KEY_ESCAPE && action == GLFW_PRESS){ + glfwSetWindowShouldClose(window, GL_TRUE); + } + else{ + switch(key){ + case GLFW_KEY_A: + renderCam->positions[0].x += (float)CAM_MOVE; + cameraReset(); + break; + case GLFW_KEY_D: + renderCam->positions[0].x -= (float)CAM_MOVE; + cameraReset(); + break; + case GLFW_KEY_W: + renderCam->positions[0].y += (float)CAM_MOVE; + cameraReset(); + break; + case GLFW_KEY_S: + renderCam->positions[0].y -= (float)CAM_MOVE; + cameraReset(); + break; + case GLFW_KEY_Q: //zoom in + renderCam->positions[0].z -= (float)CAM_MOVE; + cameraReset(); + break; + case GLFW_KEY_E: //zoom out + renderCam->positions[0].z += (float)CAM_MOVE; + cameraReset(); + break; + default: + break; + } + } +} + + +//mouse click motion function +/*void mouseClick(int button,int state, int x,int y) +{ + if(button == GLUT_RIGHT_BUTTON) + { + //std::cout<<"ss"<positions->z; + } + else if(button == GLUT_LEFT_BUTTON) + { + l_buttonDown = (state == GLUT_DOWN) ? true:false; + g_ylclick = y - (renderCam->positions->y + renderCam->views->y); + g_xlclick = x -(renderCam->positions->x + renderCam->views->x); + } +} + +void mouseMotion(int x, int y) +{ + if(r_buttonDown) + { + renderCam->positions->z = (y - g_yclick)/5.0; + iterations = 0; + finishedRender = false; + runCuda(); + //glutPostRedisplay(); + } + else if(l_buttonDown) + { + renderCam->views->x = (x - g_xlclick)/500.0; + renderCam->views->y = (y - g_ylclick)/500.0; + iterations = 0; + finishedRender = false; + runCuda(); + //glutPostRedisplay(); + } +}*/ + diff --git a/src/main.h b/src/main.h old mode 100644 new mode 100755 index 8bd2aed..67c33e6 --- a/src/main.h +++ b/src/main.h @@ -29,6 +29,7 @@ using namespace std; + //------------------------------- //----------PATHTRACER----------- //------------------------------- diff --git a/src/raytraceKernel.cu b/src/raytraceKernel.cu old mode 100644 new mode 100755 index 9c7bc7d..270cc6c --- a/src/raytraceKernel.cu +++ b/src/raytraceKernel.cu @@ -1,165 +1,535 @@ -// CIS565 CUDA Raytracer: A parallel raytracer for Patrick Cozzi's CIS565: GPU Computing at the University of Pennsylvania -// Written by Yining Karl Li, Copyright (c) 2012 University of Pennsylvania -// This file includes code from: -// Rob Farber for CUDA-GL interop, from CUDA Supercomputing For The Masses: http://www.drdobbs.com/architecture-and-design/cuda-supercomputing-for-the-masses-part/222600097 -// Peter Kutz and Yining Karl Li's GPU Pathtracer: http://gpupathtracer.blogspot.com/ -// Yining Karl Li's TAKUA Render, a massively parallel pathtracing renderer: http://www.yiningkarlli.com - -#include -#include -#include - -#include "sceneStructs.h" -#include "glm/glm.hpp" -#include "utilities.h" -#include "raytraceKernel.h" -#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); - } -} - -// 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); - - return glm::vec3((float) u01(rng), (float) u01(rng), (float) u01(rng)); -} - -// 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; -} - -//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); - - if(x<=resolution.x && y<=resolution.y){ - - 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; - } - - 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; - } -} - -// 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); - - if((x<=resolution.x && y<=resolution.y)){ - - colors[index] = generateRandomNumberFromThread(resolution, time, x, y); - } -} - -// 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))); - - // 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); - - // 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; - - // kernel launches - raytraceRay<<>>(renderCam->resolution, (float)iterations, cam, traceDepth, cudaimage, cudageoms, numberOfGeoms); - - 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 ); - delete geomList; - - // make certain the kernel has completed - cudaThreadSynchronize(); - - checkCUDAError("Kernel failed!"); -} +// CIS565 CUDA Raytracer: A parallel raytracer for Patrick Cozzi's CIS565: GPU Computing at the University of Pennsylvania +// Written by Yining Karl Li, Copyright (c) 2012 University of Pennsylvania +// This file includes code from: +// Rob Farber for CUDA-GL interop, from CUDA Supercomputing For The Masses: http://www.drdobbs.com/architecture-and-design/cuda-supercomputing-for-the-masses-part/222600097 +// Peter Kutz and Yining Karl Li's GPU Pathtracer: http://gpupathtracer.blogspot.com/ +// Yining Karl Li's TAKUA Render, a massively parallel pathtracing renderer: http://www.yiningkarlli.com + +#include +#include +#include +#include +#include "sceneStructs.h" +#include "glm/glm.hpp" +#include "utilities.h" +#include "raytraceKernel.h" +#include "intersections.h" +#include "interactions.h" + +//#define rayTracer 1 + +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); + } +} + +struct is_dead{ + __host__ __device__ bool operator()(const ray& r) + { + return r.isDead; + } +}; + +// 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); + + return glm::vec3((float) u01(rng), (float) u01(rng), (float) u01(rng)); +} + + +// 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; + glm::vec3 a = glm::normalize(glm::cross(view, up)); + glm::vec3 b = glm::normalize(glm::cross(view, a)); + glm::vec3 H = a * glm::length(view) * glm::tan(glm::radians(fov.x)); + glm::vec3 V = b * glm::length(view) * glm::tan(glm::radians(fov.y)); + glm::vec3 M = eye + view; + glm::vec3 rayDes = M + (2*((float)x/(resolution.x-1)) - 1)*H + (2*((float)y/(resolution.y-1)) - 1)*V; + //get the ray direction from eye to the destination + glm::vec3 thisRay = rayDes - eye; + + r.direction = glm::normalize(thisRay); + r.origin = eye; + r.tempColor = glm::vec3(1.0f); + r.isDead = false; + return r; +} + +//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, float iterations){ + + 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){ + + glm::vec3 color; + color.x = image[index].x*255.0f/iterations; + color.y = image[index].y*255.0f/iterations; + color.z = image[index].z*255.0f/iterations; //weight for each iteration + + 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; + } +} + + +// loop through all geometry to test ray intersection, returns the geoID that corresponds to intersected geometry +__host__ __device__ int intersectTest(ray r, glm::vec3& intersect, glm::vec3& normal, staticGeom* geoms, int numberOfGeoms, triangle * cudatris){ + + if(r.isDead) + return -1; //didn't hit anything + float distMin = -2, dist = -1; + glm::vec3 tempIntersect(0.0f); + glm::vec3 tempNormal(0.0f); + int ID = -1; + + for (int g=0; g -0.5f ) || ( distMin > -1 && dist < distMin && dist > -0.5f ) ){ + distMin = dist; + ID = g; + intersect = tempIntersect; + normal = tempNormal; + } + } + return ID; +} + +//return true if ray directly hit lights +__host__ __device__ bool LightRayTest(ray r, staticGeom* geoms, int numberOfGeoms, material* materials, triangle * cudatris){ + glm::vec3 intersPoint(0.0f); + glm::vec3 intersNormal(0.0f); + + //printf("shadow ray: [%f,%f,%f], [%f,%f,%f]\n", sr.origin.x,sr.origin.y,sr.origin.z,sr.direction.x,sr.direction.y,sr.direction.z); + int geoID = intersectTest(r, intersPoint, intersNormal, geoms, numberOfGeoms, cudatris); + if( geoID>-1 && materials[geoms[geoID].materialid].emittance > 0){ //hit light soource + return true; + } + else{ + return false; + } + +} + +//calculates the direct lighting for a certain hit point and modify color of that hit +__device__ __host__ void directLighting(float seed, glm::vec3& theColor, glm::vec3& theIntersect, glm::vec3& theNormal, int geoID, int* lights, int numOfLights, material* cudamats, staticGeom* geoms, int numOfGeoms, triangle * cudatris){ + ray shadowRay; + float rayLen; + float lightArea; + glm::vec3 lightNormal; + + int chosenLight = lights[0]; + if( numOfLights > 1){ + thrust::default_random_engine rng(hash(seed)); + thrust::uniform_real_distribution u01(0,1); + chosenLight = lights[(int)((float)u01(rng) * numOfLights)]; + } + glm::vec3 Plight; + if( geoms[chosenLight].type == CUBE ){ + Plight = getRandomPointOnCube( geoms[chosenLight], seed); + } + else if( geoms[chosenLight].type == SPHERE ){ + Plight = getRandomPointOnSphere( geoms[chosenLight], seed); + } + + shadowRay.direction = glm::normalize(Plight - theIntersect); + shadowRay.origin = theIntersect + (float)EPSILON * shadowRay.direction; + int lightID = glm::length(Plight - theIntersect); + + material curMat = cudamats[geoms[geoID].materialid]; //material of the hit goemetry + if(LightRayTest(shadowRay, geoms, numOfGeoms, cudamats, cudatris)){ + float cosTerm = glm::clamp( glm::dot( theNormal, shadowRay.direction ), 0.0f, 1.0f); //proportion of facing light + float cosTerm2 = glm::clamp( glm::dot( lightNormal, -shadowRay.direction ), 0.0f, 1.0f); //proportion of incoming light + float areaSampling = lightArea / (float) pow( rayLen, 2.0f) ; // dA/r^2 + theColor += cudamats[lightID].emittance * curMat.color * cosTerm * cosTerm2 * areaSampling ; + } +} + +#ifdef rayTracer +//TODO: IMPLEMENT THIS FUNCTION +//Core raytracer kernel (recursive) + +__host__ __device__ glm::vec3 raytraceRecursive(ray r, int iteration, float currentIndexOfRefraction, int depth, int maximumDepth, staticGeom* geoms, int numberOfGeoms, material* materials, int numberOfMaterials, light* lightSources, int numberOfLights){ + + glm::vec3 bgColor(0.0f); + glm::vec3 ambientColor(1.0f); + glm::vec3 phongColor(0.0f), reflColor(0.0f), refraColor(0.0f);; + glm::vec3 returnColor(0.0f); + float ka = 0.2f; + + if(depth > maximumDepth) + return bgColor; + + // intersection test + glm::vec3 intersectionPoint, intersectionNormal; + int intersIndex = rayIntersect(r, geoms, numberOfGeoms, intersectionPoint, intersectionNormal, materials); + material mat = materials[geoms[intersIndex].materialid]; + + if(intersIndex == -1) return bgColor; + + else if(mat.emittance > 0.0f){ // intersected with light source geometry + returnColor = mat.color; + } + + else{ // intersected with actual geometry + +// returnColor = ka * ambientColor * materials[geoms[intersIndex].materialid].color; + + if(/*iteration == 0 && */materials[geoms[intersIndex].materialid].hasRefractive == 1) + { + float nextIndexOfRefraction = 1.0f; + glm::vec3 refraDir; + if(abs(currentIndexOfRefraction - 1) < 0.00001f) // current ray is in air + { + refraDir = calculateRefractionDirection(r.direction, intersectionNormal, currentIndexOfRefraction, materials[geoms[intersIndex].materialid].indexOfRefraction, nextIndexOfRefraction); + } + else // current ray is in glass + { + refraDir = calculateRefractionDirection(r.direction, -intersectionNormal, currentIndexOfRefraction, 1.0f, nextIndexOfRefraction); + } + + ray refraRay; + refraRay.origin = intersectionPoint + 0.01f * refraDir; + refraRay.direction = refraDir; + refraColor = raytraceRecursive(refraRay, iteration, nextIndexOfRefraction, depth + 1, maximumDepth, geoms, numberOfGeoms, materials, numberOfMaterials, lightSources, numberOfLights); + returnColor += refraColor; + } + + if(materials[geoms[intersIndex].materialid].hasReflective == 1) + { + glm::vec3 reflDir = calculateReflectionDirection(intersectionNormal, r.direction); + ray reflRay; + reflRay.origin = intersectionPoint + 0.01f * reflDir; + reflRay.direction = reflDir; + reflColor = raytraceRecursive(reflRay, iteration, 1.0f, depth + 1, maximumDepth, geoms, numberOfGeoms, materials, numberOfMaterials, lightSources, numberOfLights); + returnColor += reflColor; + } + + + if(iteration < numberOfLights){ + if(ShadowRayUnblocked(intersectionPoint, lightSources[iteration].position, geoms, numberOfGeoms, materials)) + { + glm::vec3 L = glm::normalize(lightSources[iteration].position - intersectionPoint); + float dot1 = glm::clamp(glm::dot(intersectionNormal, L), 0.0f, 1.0f); + float dot2 = glm::dot(calculateReflectionDirection(intersectionNormal, -L) ,-r.direction); + glm::vec3 diffuse = lightSources[iteration].color * 0.5f * materials[geoms[intersIndex].materialid].color * dot1; + glm::vec3 specular; + if(abs(materials[geoms[intersIndex].materialid].specularExponent) > 1e-6) + specular = lightSources[iteration].color * 0.1f * pow(glm::max(dot2, 0.0f), materials[geoms[intersIndex].materialid].specularExponent); + phongColor += diffuse + specular; + + } + } + + returnColor += (5.0f / numberOfLights) * (0.1f * (float)numberOfLights * reflColor + (float)numberOfLights * refraColor); + } + return returnColor; +} + + +__global__ void raytracePrimary(glm::vec2 resolution, int time, cameraData cam, int rayDepth, glm::vec3* colors, + staticGeom* geoms, int numberOfGeoms, material* materials, int numberOfMaterials, int* lightSources, int numberOfLights){ + + int x = (blockIdx.x * blockDim.x) + threadIdx.x; + int y = (blockIdx.y * blockDim.y) + threadIdx.y; + int index = x + (y * resolution.x); + + //on every thread, get color for any pixel given pixel(x,y) and camera + if((x<=resolution.x && y<=resolution.y)){ + int init_depth = 0; + ray r = raycastFromCameraKernel(resolution, time, x, y, cam.position, cam.view, cam.up, cam.fov); + colors[index] += raytraceRecursive(r, time, 1.0f, init_depth, rayDepth, geoms, numberOfGeoms, materials, numberOfMaterials, lightSources, numberOfLights); + } + +} +#endif + +// TODO: IMPLEMENT THIS FUNCTION +// Core path tracer kernel +__global__ void pathtraceRay(ray* rays, float time, int rayDepth, int numOfRays, glm::vec3* colors, staticGeom* geoms, int numberOfGeoms, material* cudamats, int* lights, int numOfLights, cameraData cam, triangle* cudatris){ + + int x = (blockIdx.x * blockDim.x) + threadIdx.x; + int y = (blockIdx.y * blockDim.y) + threadIdx.y; + int index = x + (int)ceil(sqrt((float)numOfRays))* y; + + if( index < numOfRays ){ + float seed = (float)index * (float)time * ( (float)rayDepth + 1.0f ); + ray r = rays[index]; + + glm::vec3 Pintersect(0.0f); + glm::vec3 Pnormal(0.0f); + int intersIndex = intersectTest(r, Pintersect, Pnormal, geoms, numberOfGeoms, cudatris); + + if(intersIndex!=-1){ + material curMat = cudamats[geoms[intersIndex].materialid]; + if( curMat.emittance > 0 ){ //ray ends when hit light source + colors[r.pixelIndex] += r.tempColor * curMat.color * curMat.emittance; + r.isDead = true; + } + else{ // for reflection and refraction effect + if(curMat.hasReflective > 0 || curMat.hasRefractive > 0){ + Fresnel Fres; + float reflectance; + glm::vec3 reflectDir, transmitDir; + if(glm::dot(r.direction,Pnormal)<0){ //ray is outside + Fres = calculateFresnel(Pnormal,r.direction,1.0f, curMat.indexOfRefraction); + reflectDir = calculateReflectionDirection(Pnormal, r.direction); + transmitDir = calculateTransmissionDirection(Pnormal, r.direction, 1.0f, curMat.indexOfRefraction); + } + else{ //ray is inside + Fres = calculateFresnel(-Pnormal,r.direction, curMat.indexOfRefraction, 1.0f); + reflectDir = calculateReflectionDirection(-Pnormal, r.direction); + transmitDir = calculateTransmissionDirection(-Pnormal, r.direction, curMat.indexOfRefraction, 1.0f); + } + + if( curMat.hasRefractive > 0 && curMat.hasReflective > 0){ + thrust::default_random_engine rng( hash( seed ) ); + thrust::uniform_real_distribution u01(0,1); + + if((float) u01(rng) < Fres.reflectionCoefficient ){ //reflected + r.direction = reflectDir; + } + else{ //transmitted + r.direction = transmitDir; + } + } + else if(curMat.hasReflective > 0){ + r.direction = reflectDir; + } + else if (curMat.hasRefractive > 0){ + r.direction = transmitDir; + } + r.origin = Pintersect + (float)EPSILON * r.direction; + if(glm::length(curMat.color)>0) + r.tempColor *= curMat.color ; + } + + else{ + thrust::default_random_engine rng(hash(seed)); + thrust::uniform_real_distribution u01(0,1); + if((float) u01(rng) < 0.01f ){ //direct light + directLighting(seed,r.tempColor,Pintersect,Pnormal,intersIndex,lights,numOfLights, cudamats,geoms, numberOfGeoms, cudatris); + } + else{ + //cos weighted + r.direction = calculateCosWeightedRandomDirInHemisphere(Pnormal, (float) u01(rng), (float) u01(rng)); + r.origin = Pintersect + (float)EPSILON * r.direction ; + float diffuseTerm = glm::clamp( glm::dot( Pnormal,r.direction ), 0.0f, 1.0f); + r.tempColor *= diffuseTerm * curMat.color; + } + } + } + } + else{ //if ray hit nothing + r.isDead = true; + } + rays[index] = r; + } +} + +//initialize the ray pool for cudarays +__global__ void generateRaypool(ray * rayPool, cameraData cam, float iterations,glm::vec3 *colors, staticGeom* geoms, int numberOfGeoms, material* cudamats, int * lightIDs, int numberOfLights, triangle * cudatris){ + int x = (blockIdx.x * blockDim.x) + threadIdx.x; + int y = (blockIdx.y * blockDim.y) + threadIdx.y; + int index = x + (y * cam.resolution.x); + //ray r = rayPool[index]; + + if( x<= cam.resolution.x && y <= cam.resolution.y ){ + ray r = raycastFromCameraKernel( cam.resolution, iterations, x, y, cam.position, cam.view, cam.up, cam.fov ); + r.pixelIndex = index; + + if(DEPTH_OF_FIELD){ + glm::vec3 focalPoint = r.origin + r.direction * cam.focalLength / glm::dot(cam.view, r.direction); //L = f/cos(theta) + thrust::default_random_engine rng(hash((float)index*iterations)); + thrust::uniform_real_distribution u01(0,1); + float theta = 2.0f * PI * u01(rng); + float radius = u01(rng) * cam.aperture; + glm::vec3 eyeOffset(cos(theta)*radius, sin(theta)*radius, 0); + glm::vec3 newEyePoint = cam.position + eyeOffset; + r.origin = newEyePoint; + r.direction = glm::normalize(focalPoint - newEyePoint); + } + + glm::vec3 Pintersect(0.0f); + glm::vec3 Pnormal(0.0f); + int geoID = intersectTest(r, Pintersect, Pnormal, geoms, numberOfGeoms, cudatris); + if( geoID > -1){ + directLighting((float)index*iterations, colors[index], Pintersect, Pnormal,geoID, lightIDs, numberOfLights, cudamats, geoms, numberOfGeoms, cudatris); + } + rayPool[index] = r; + } +} + + +// TODO: FINISH THIS FUNCTION +// Wrapper for the __global__ call that sets up the kernel calls and does a ton of memory management +void cudaPathTraceCore(uchar4* PBOpos, camera* renderCam, int frame, int iterations, material* materials, int numberOfMaterials, geom* geoms, int numberOfGeoms){ + + int traceDepth = 10; //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))); + + // 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); + + // package geometry and materials and sent to GPU + staticGeom* geomList = new staticGeom[numberOfGeoms]; + int meshID = -1; + triangle* cudatris = NULL; + + for(int i=0; iresolution; + cam.position = renderCam->positions[frame]; + cam.view = renderCam->views[frame]; + cam.up = renderCam->ups[frame]; + cam.fov = renderCam->fov; + cam.aperture = renderCam->aperture; + cam.focalLength = renderCam->focalLength; + + // material setup + material* cudamats = NULL; + cudaMalloc((void**)&cudamats, numberOfMaterials*sizeof(material)); + cudaMemcpy( cudamats, materials, numberOfMaterials*sizeof(material), cudaMemcpyHostToDevice); + + //lights setup + int numberOfLights = 0; + for(int i = 0; i < numberOfGeoms; ++i){ + if(materials[geoms[i].materialid].emittance > 0){ + numberOfLights ++ ; + } + } + + int *lightIDs = new int[numberOfLights]; + int k = 0; + for(int i = 0; i < numberOfGeoms; ++i){ + if(materials[geoms[i].materialid].emittance > 0){ + lightIDs[k] = i; + k++; + } + } + int* cudalightIDs = NULL; + cudaMalloc((void**)&cudalightIDs, numberOfLights*sizeof(int)); + cudaMemcpy( cudalightIDs, lightIDs, numberOfLights*sizeof(int), cudaMemcpyHostToDevice); + + + //set up ray pool on device + ray* cudarays = NULL; + int numOfRays = cam.resolution.x * cam.resolution.y; + cudaMalloc((void**)&cudarays, numOfRays*sizeof(ray)); + generateRaypool<<>>(cudarays, cam, (float)iterations, cudaimage, cudageoms, numberOfGeoms, cudamats, cudalightIDs, numberOfLights, cudatris); + + for(int cur_depth=0; cur_depth0; cur_depth++){ + + thrust::device_ptr raypoolStart = thrust::device_pointer_cast(cudarays); //coverts cuda pointer to thrust pointer + thrust::device_ptr raypoolEnd = thrust::remove_if(raypoolStart, raypoolStart + numOfRays, is_dead()); + numOfRays = (int)(raypoolEnd-raypoolStart); + + //xBlocks * yBlocks = numOfRays / (tileSize*tileSize) + int xBlocks = (int) ceil( sqrt((float)numOfRays)/(float)(tileSize) ); + int yBlocks = (int) ceil( sqrt((float)numOfRays)/(float)(tileSize) ); + dim3 newBlocksPerGrid(xBlocks,yBlocks); + + pathtraceRay<<>>(cudarays, (float)iterations, cur_depth, (int)numOfRays, cudaimage, cudageoms, numberOfGeoms, cudamats, cudalightIDs, numberOfLights, cam, cudatris); + } + + //raytraceRay<<>>(cudarays, (float)iterations, cur_depth, (int)numOfRays, cudaimage, cudageoms, numberOfGeoms, cudamats, cudalightIDs, numberOfLights, cam, cudatris); + sendImageToPBO<<>>(PBOpos, renderCam->resolution, cudaimage,(float)iterations); + + // 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( cudamats ); + cudaFree( cudarays ); + cudaFree( cudalightIDs ); + if(meshID>-1){ + cudaFree( cudatris ); + } + + delete geomList; + delete lightIDs; + // make certain the kernel has completed + cudaThreadSynchronize(); + + checkCUDAError("Kernel failed!"); +} diff --git a/src/raytraceKernel.h b/src/raytraceKernel.h old mode 100644 new mode 100755 index e527a81..58608da --- a/src/raytraceKernel.h +++ b/src/raytraceKernel.h @@ -5,15 +5,20 @@ // Peter Kutz and Yining Karl Li's GPU Pathtracer: http://gpupathtracer.blogspot.com/ // Yining Karl Li's TAKUA Render, a massively parallel pathtracing renderer: http://www.yiningkarlli.com -#ifndef RAYTRACEKERNEL_H -#define PATHTRACEKERNEL_H #include -#include #include #include +#include +#include +#include +#include +#include #include "sceneStructs.h" -void cudaRaytraceCore(uchar4* pos, camera* renderCam, int frame, int iterations, material* materials, int numberOfMaterials, geom* geoms, int numberOfGeoms); +#define THRESHOLD 0.001 +#define DEPTH_OF_FIELD 0 + +void cudaRayTraceCore(uchar4* pos, camera* renderCam, int frame, int iterations, material* materials, int numberOfMaterials, geom* geoms, int numberOfGeoms); +void cudaPathTraceCore(uchar4* pos, camera* renderCam, int frame, int iterations, material* materials, int numberOfMaterials, geom* geoms, int numberOfGeoms); -#endif diff --git a/src/scene.cpp b/src/scene.cpp old mode 100644 new mode 100755 index 4cbe216..817c409 --- a/src/scene.cpp +++ b/src/scene.cpp @@ -54,20 +54,8 @@ int scene::loadObject(string objectid){ cout << "Creating new cube..." << endl; newObject.type = CUBE; }else{ - string objline = line; - string name; - string extension; - istringstream liness(objline); - getline(liness, name, '.'); - getline(liness, extension, '.'); - if(strcmp(extension.c_str(), "obj")==0){ - cout << "Creating new mesh..." << endl; - cout << "Reading mesh from " << line << "... " << endl; - newObject.type = MESH; - }else{ - cout << "ERROR: " << line << " is not a valid object type!" << endl; - return -1; - } + cout<< "error object type" < tokens = utilityCore::tokenizeString(line); @@ -152,6 +140,10 @@ int scene::loadCamera(){ newCamera.iterations = atoi(tokens[1].c_str()); }else if(strcmp(tokens[0].c_str(), "FILE")==0){ newCamera.imageName = tokens[1]; + }else if(strcmp(tokens[0].c_str(), "FOCAL")==0){ + newCamera.focalLength = atof(tokens[1].c_str()); + }else if(strcmp(tokens[0].c_str(), "APERTURE")==0){ + newCamera.aperture = atof(tokens[1].c_str()); } } diff --git a/src/scene.h b/src/scene.h old mode 100644 new mode 100755 diff --git a/src/sceneStructs.h b/src/sceneStructs.h old mode 100644 new mode 100755 index 5e0c853..edf3dba --- a/src/sceneStructs.h +++ b/src/sceneStructs.h @@ -6,16 +6,28 @@ #ifndef CUDASTRUCTS_H #define CUDASTRUCTS_H +#include #include "glm/glm.hpp" #include "cudaMat4.h" #include #include + enum GEOMTYPE{ SPHERE, CUBE, MESH }; struct ray { + + glm::vec3 origin; glm::vec3 direction; + + bool isDead; + int pixelIndex; + glm::vec3 tempColor; +}; + +struct triangle{ + glm::vec3 p1,p2,p3; }; struct geom { @@ -27,6 +39,11 @@ struct geom { glm::vec3* scales; cudaMat4* transforms; cudaMat4* inverseTransforms; + + glm::vec3 boundingBoxMin; + glm::vec3 boundingBoxMax; + triangle* tris; + int numOfTris; }; struct staticGeom { @@ -37,6 +54,11 @@ struct staticGeom { glm::vec3 scale; cudaMat4 transform; cudaMat4 inverseTransform; + + glm::vec3 boundingBoxMin; + glm::vec3 boundingBoxMax; + triangle* tris; + int numOfTris; }; struct cameraData { @@ -45,6 +67,9 @@ struct cameraData { glm::vec3 view; glm::vec3 up; glm::vec2 fov; + + float aperture; + float focalLength; }; struct camera { @@ -58,6 +83,9 @@ struct camera { glm::vec3* image; ray* rayList; std::string imageName; + + float focalLength; + float aperture; }; struct material{ @@ -73,4 +101,6 @@ struct material{ float emittance; }; + + #endif //CUDASTRUCTS_H diff --git a/src/utilities.cpp b/src/utilities.cpp old mode 100644 new mode 100755 index 1abf9b6..0c579f2 --- a/src/utilities.cpp +++ b/src/utilities.cpp @@ -4,12 +4,15 @@ // File: utilities.cpp // A collection/kitchen sink of generally useful functions +//#define GLM_FORCE_RADIANS + #include #include #include #include "utilities.h" + float utilityCore::clamp(float f, float min, float max){ if(f - - - - Debug - Win32 - - - Release - Win32 - - - - {B7AAC719-5B66-4370-9FD7-8ED6DB66347B} - Project3Pathtracer - - - - Application - true - MultiByte - - - Application - false - true - MultiByte - - - - - - - - - - - - - - - - Level3 - Disabled - $(SolutionDir)/../../external/include/;%(AdditionalIncludeDirectories) - - - true - $(SolutionDir)/../../external/lib/win/GLFW/;$(SolutionDir)/../../external/lib/win/GL/;%(AdditionalLibraryDirectories) - cudart.lib;glew32s.lib;glfw3.lib;opengl32.lib;%(AdditionalDependencies) - Console - - - - - Level3 - MaxSpeed - true - true - $(SolutionDir)/../../external/include/;%(AdditionalIncludeDirectories) - - - true - true - true - $(SolutionDir)/../../external/lib/win/GLFW/;$(SolutionDir)/../../external/lib/win/GL/;%(AdditionalLibraryDirectories) - cudart.lib;glew32s.lib;glfw3.lib;opengl32.lib;%(AdditionalDependencies) - Console - - - - - - - - - - - - - - - - - - - - - - - - - Document - - - - - - + + + + + Debug + Win32 + + + Release + Win32 + + + + {B7AAC719-5B66-4370-9FD7-8ED6DB66347B} + Project3Pathtracer + 565-PathTracing + + + + Application + true + MultiByte + + + Application + false + true + MultiByte + + + + + + + + + + + + + + + + Level3 + Disabled + $(SolutionDir)/../../external/include/; Q:\CUDA5.5\include;Q:\CUDA5.5\common\inc; %(AdditionalIncludeDirectories) + + + true + $(SolutionDir)/../../external/lib/win/GLFW/;$(SolutionDir)/../../external/lib/win/GL/;%(AdditionalLibraryDirectories) + cudart.lib;glew32s.lib;glfw3.lib;opengl32.lib;%(AdditionalDependencies) + Console + + + compute_30,sm_30 + + + + + Level3 + MaxSpeed + true + true + $(SolutionDir)/../../external/include/;%(AdditionalIncludeDirectories) + + + true + true + true + $(SolutionDir)/../../external/lib/win/GLFW/;$(SolutionDir)/../../external/lib/win/GL/;%(AdditionalLibraryDirectories) + cudart.lib;glew32s.lib;glfw3.lib;opengl32.lib;%(AdditionalDependencies) + Console + + + compute_30,sm_30 + $(SolutionDir)/../../external/include/; + + + + + + + + + + + + + + + + + + + + + + + + + Document + + + + + + + + + \ 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 old mode 100644 new mode 100755 index 584fd19..d4c1bc8 --- a/windows/Project3-Pathtracer/Project3-Pathtracer/Project3-Pathtracer.vcxproj.filters +++ b/windows/Project3-Pathtracer/Project3-Pathtracer/Project3-Pathtracer.vcxproj.filters @@ -1,80 +1,85 @@ - - - - - {4FC737F1-C7A5-4376-A066-2A32D752A2FF} - cpp;c;cc;cxx;def;odl;idl;hpj;bat;asm;asmx - - - {93995380-89BD-4b04-88EB-625FBE52EBFB} - h;hpp;hxx;hm;inl;inc;xsd - - - {67DA6AB6-F800-4c08-8B7A-83BB121AAD01} - rc;ico;cur;bmp;dlg;rc2;rct;bin;rgs;gif;jpg;jpeg;jpe;resx;tiff;tif;png;wav;mfcribbon-ms - - - {2ff2df1e-6ae6-4823-885d-b970fbc86988} - - - {f1a3a052-d6b3-4aff-a9f3-ac85b6c99146} - - - - - Header Files - - - Header Files - - - Header Files - - - Header Files - - - Header Files - - - Header Files - - - Header Files - - - Header Files - - - Header Files - - - - - Source Files - - - Source Files - - - Source Files - - - Source Files - - - Source Files\glslUtil - - - Source Files\stb_image - - - Source Files\stb_image - - - - - Source Files - - + + + + + {4FC737F1-C7A5-4376-A066-2A32D752A2FF} + cpp;c;cc;cxx;def;odl;idl;hpj;bat;asm;asmx + + + {93995380-89BD-4b04-88EB-625FBE52EBFB} + h;hpp;hxx;hm;inl;inc;xsd + + + {67DA6AB6-F800-4c08-8B7A-83BB121AAD01} + rc;ico;cur;bmp;dlg;rc2;rct;bin;rgs;gif;jpg;jpeg;jpe;resx;tiff;tif;png;wav;mfcribbon-ms + + + {2ff2df1e-6ae6-4823-885d-b970fbc86988} + + + {f1a3a052-d6b3-4aff-a9f3-ac85b6c99146} + + + + + Header Files + + + Header Files + + + Header Files + + + Header Files + + + Header Files + + + Header Files + + + Header Files + + + Header Files + + + Header Files + + + + + Source Files + + + Source Files + + + Source Files + + + Source Files + + + Source Files\glslUtil + + + Source Files\stb_image + + + Source Files\stb_image + + + + + Source Files + + + + + Resource Files + + \ No newline at end of file