diff --git a/README.md b/README.md index 6bef2b9..02ee087 100644 --- a/README.md +++ b/README.md @@ -3,331 +3,77 @@ CIS565: Project 1: CUDA Raytracer ------------------------------------------------------------------------------- Fall 2013 ------------------------------------------------------------------------------- -Due Thursday, 09/19/2013 -------------------------------------------------------------------------------- - -------------------------------------------------------------------------------- -NOTE: -------------------------------------------------------------------------------- -This project requires an NVIDIA graphics card with CUDA capability! Any card -after the Geforce 8xxx series will work. If you do not have an NVIDIA graphics -card in the machine you are working on, feel free to use any machine in the SIG -Lab or in Moore100 labs. All machines in the SIG Lab and Moore100 are equipped -with CUDA capable NVIDIA graphics cards. If this too proves to be a problem, -please contact Patrick or Liam as soon as possible. ------------------------------------------------------------------------------- -INTRODUCTION: +IMPLEMENTATION: ------------------------------------------------------------------------------- -In this project, you will implement a CUDA based raytracer capable of -generating raytraced rendered images extremely quickly. For those of you who -have taken CIS460/560, building a raytracer should not be anything new to you -from a conceptual point of you. For those of you that have not taken -CIS460/560, raytracing is a technique for generating images by tracing rays of -light through pixels in an image plane out into a scene and following the way -the rays of light bounce and interact with objects in the scene. More -information can be found here: -http://en.wikipedia.org/wiki/Ray_tracing_(graphics). - -The ultimate purpose of this project is to serve as the foundation for your -next project: a full CUDA based global illumination pathtracer. Raytracing can -be thought of as a way to generate an isolated version of the direct light -contribution in a global illumination scenario. - -Since in this class we are concerned with working in generating actual images -and less so with mundane tasks like file I/O, this project includes basecode -for loading a scene description file format, described below, and various other -things that generally make up the render "harness" that takes care of -everything up to the rendering itself. The core renderer is left for you to -implement. Finally, note that while this basecode is meant to serve as a -strong starting point for a CUDA raytracer, you are not required to use this -basecode if you wish, and you may also change any part of the basecode -specification as you please, so long as the final rendered result is correct. - -------------------------------------------------------------------------------- -CONTENTS: -------------------------------------------------------------------------------- -The Project1 root directory contains the following subdirectories: - -* src/ contains the source code for the project. Both the Windows Visual Studio - solution and the OSX and Linux makefiles reference this folder for all - source; the base source code compiles on Linux, OSX and Windows without - modification. -* scenes/ contains an example scene description file. -* renders/ contains an example render of the given example scene file. -* PROJ1_WIN/ contains a Windows Visual Studio 2010 project and all dependencies - needed for building and running on Windows 7. -* PROJ1_OSX/ contains a OSX makefile, run script, and all dependencies needed - for building and running on Mac OSX 10.8. -* PROJ1_NIX/ contains a Linux makefile for building and running on Ubuntu - 12.04 LTS. Note that you will need to set the following environment - variables: - - - PATH=$PATH:/usr/local/cuda-5.5/bin - - LD_LIBRARY_PATH=/usr/local/cuda-5.5/lib64:/lib - - you may set these any way that you like. I added them to my .bashrc - - -The Windows and OSX versions of the project build and run exactly the same way -as in Project0. +For this project, I learned about CUDA and implemented a simple GPU raytracer. +The basic features include diffuse and phong shading of cubes and spheres with +raytraced shadows from a single point light. I also added these additional +features: -------------------------------------------------------------------------------- -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 -* Phong lighting for one point light source -* Diffuse lambertian surfaces -* Raytraced shadows -* Cube intersection testing -* Sphere surface point sampling - -You are also required to implement at least 2 of the following features: - -* Specular reflection * Soft shadows and area lights -* Texture mapping -* Bump mapping * Depth of field -* Supersampled antialiasing -* Refraction, i.e. glass -* OBJ Mesh loading and renderin -* Interactive camera - -------------------------------------------------------------------------------- -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. +* Supersampled anti-aliasing +* Movable camera -* raytraceKernel.cu contains the core raytracing CUDA kernel. You will need to - complete: - * cudaRaytraceCore() handles kernel launches and memory management; this - function already contains example code for launching kernels, - transferring geometry and cameras from the host to the device, and transferring - image buffers from the host to the device and back. You will have to complete - this function to support passing materials and lights to CUDA. - * raycastFromCameraKernel() is a function that you need to implement. This - function once correctly implemented should handle camera raycasting. - * raytraceRay() is the core raytracing CUDA kernel; all of your raytracing - logic should be implemented in this CUDA kernel. raytraceRay() should - 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. +Please use the new scene files provided in the repository. phongMatScene.txt is +the scene I used to render the images below. +The scenes are still read through the command line. I added focal length and +aperture as new camera features that can be read in through the scene files, +so the old scene files would not work. -* 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. +The camera controls are: +* 'a' and 'd' keys to translate along the x axis +* 'w' and 's' keys to translate along the y axis +* 'j' and 'k' keys to translate along the z axis -* 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. ------------------------------------------------------------------------------- -TAKUAscene FORMAT: +SCREENSHOT OF MY RAYTRACER ------------------------------------------------------------------------------- -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 +Screen shot +![Alt text](/renders/screenCap/sampleSceneDOF.jpg "screen shot") -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 +Depth of Field +![Alt text](/renders/screenCap/DOF.bmp "DOF") -An example TAKUAscene file setting up two frames inside of a Cornell Box can be -found in the scenes/ directory. +No DOF +![Alt text](/renders/screenCap/noDOF.bmp "No DOF") -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. +Two lights with depth of Field +![Alt text](/renders/screenCap/twoLightDOF.bmp "two lights with DOF") ------------------------------------------------------------------------------- -README +VIDEO ------------------------------------------------------------------------------- -All students must replace or augment the contents of this Readme.md in a clear -manner with the following: +Video link: -* A brief description of the project and the specific features you implemented. -* At least one screenshot of your project running. -* A 30 second or longer video of your project running. To create the video you - can use http://www.microsoft.com/expression/products/Encoder4_Overview.aspx -* A performance evaluation (described in detail below). +https://vimeo.com/75069384 ------------------------------------------------------------------------------- PERFORMANCE EVALUATION ------------------------------------------------------------------------------- -The performance evaluation is where you will investigate how to make your CUDA -programs more efficient using the skills you've learned in class. You must have -perform at least one experiment on your code to investigate the positive or -negative effects on performance. - -One such experiment would be to investigate the performance increase involved -with adding a spatial data-structure to your scene data. - -Another idea could be looking at the change in timing between various block -sizes. +![Alt text](/renders/performanceAnalysis.png "Performance Analysis") -A good metric to track would be number of rays per second, or frames per -second, or number of objects displayable at 60fps. +Link to pdf -We encourage you to get creative with your tweaks. Consider places in your code -that could be considered bottlenecks and try to improve them. - -Each student should provide no more than a one page summary of their -optimizations along with tables and or graphs to visually explain and -performance differences. +https://github.com/zxyzhu/Project1-RayTracer/blob/master/renders/PerformanceAnalysis.pdf ------------------------------------------------------------------------------- -THIRD PARTY CODE POLICY +THIRD PARTY CODE ------------------------------------------------------------------------------- -* 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. +* None, but here are the resources I used for box intersection and sampling +random points on spheres: -------------------------------------------------------------------------------- -SELF-GRADING -------------------------------------------------------------------------------- -* On the submission date, email your grade, on a scale of 0 to 100, to Liam, - liamboone+cis565@gmail.com, with a one paragraph explanation. Be concise and - realistic. Recall that we reserve 30 points as a sanity check to adjust your - grade. Your actual grade will be (0.7 * your grade) + (0.3 * our grade). We - hope to only use this in extreme cases when your grade does not realistically - reflect your work - it is either too high or too low. In most cases, we plan - to give you the exact grade you suggest. -* Projects are not weighted evenly, e.g., Project 0 doesn't count as much as - the path tracer. We will determine the weighting at the end of the semester - based on the size of each project. +* Box intersection: +http://www.siggraph.org/education/materials/HyperGraph/raytrace/rtinter3.htm + +* Sphere random point generation: +http://mathworld.wolfram.com/SpherePointPicking.html ------------------------------------------------------------------------------- -SUBMISSION +FUTURE WORK ------------------------------------------------------------------------------- -As with the previous project, you should fork this project and work inside of -your fork. Upon completion, commit your finished project back to your fork, and -make a pull request to the master repository. You should include a README.md -file in the root directory detailing the following - -* A brief description of the project and specific features you implemented -* At least one screenshot of your project running, and at least one screenshot - of the final rendered output of your raytracer -* A link to a video of your raytracer running. -* Instructions for building and running your project if they differ from the - base code -* A performance writeup as detailed above. -* A list of all third-party code used. -* This Readme file, augmented or replaced as described above in the README section. +I am currently working on reflective surfaces. After that I plan to make +refraction and create better camera controls. diff --git a/renders/PerformanceAnalysis.pdf b/renders/PerformanceAnalysis.pdf new file mode 100644 index 0000000..e60834d Binary files /dev/null and b/renders/PerformanceAnalysis.pdf differ diff --git a/renders/performanceAnalysis.png b/renders/performanceAnalysis.png new file mode 100644 index 0000000..edf8831 Binary files /dev/null and b/renders/performanceAnalysis.png differ diff --git a/renders/screenCap/DOF.bmp b/renders/screenCap/DOF.bmp new file mode 100644 index 0000000..b68e19d Binary files /dev/null and b/renders/screenCap/DOF.bmp differ diff --git a/renders/screenCap/noDOF.bmp b/renders/screenCap/noDOF.bmp new file mode 100644 index 0000000..d044107 Binary files /dev/null and b/renders/screenCap/noDOF.bmp differ diff --git a/renders/screenCap/sampleSceneDOF.jpg b/renders/screenCap/sampleSceneDOF.jpg new file mode 100644 index 0000000..395b9a2 Binary files /dev/null and b/renders/screenCap/sampleSceneDOF.jpg differ diff --git a/renders/screenCap/twoLightDOF.bmp b/renders/screenCap/twoLightDOF.bmp new file mode 100644 index 0000000..c4bcb29 Binary files /dev/null and b/renders/screenCap/twoLightDOF.bmp differ diff --git a/renders/screenCap/twoLights.bmp b/renders/screenCap/twoLights.bmp new file mode 100644 index 0000000..2bcc115 Binary files /dev/null and b/renders/screenCap/twoLights.bmp differ diff --git a/scenes/phongMatScene.txt b/scenes/phongMatScene.txt new file mode 100644 index 0000000..381626f --- /dev/null +++ b/scenes/phongMatScene.txt @@ -0,0 +1,191 @@ +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 50 +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 20 +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 25 +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 + +CAMERA +RES 800 800 +FOVY 25 +ITERATIONS 5000 +FILE test.bmp +frame 0 +EYE 0 4.5 15 +VIEW 0 0 -1 +UP 0 1 0 +FOCAL 12 +APERTURE 0.3 + +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 +sphere +material 4 +frame 0 +TRANS 0 2 0 +ROTAT 0 180 0 +SCALE 3 3 3 + +OBJECT 6 +sphere +material 3 +frame 0 +TRANS 2 5 2 +ROTAT 0 180 0 +SCALE 2.5 2.5 2.5 + +OBJECT 7 +sphere +material 6 +frame 0 +TRANS -2 5 -2 +ROTAT 0 180 0 +SCALE 3 3 3 + +OBJECT 8 +cube +material 8 +frame 0 +TRANS 0 10 0 +ROTAT 0 0 90 +SCALE .3 3 3 \ No newline at end of file diff --git a/scenes/sampleScene.txt b/scenes/sampleScene.txt index 877706f..2cb7c67 100755 --- a/scenes/sampleScene.txt +++ b/scenes/sampleScene.txt @@ -115,6 +115,8 @@ frame 0 EYE 0 4.5 12 VIEW 0 0 -1 UP 0 1 0 +FOCAL 12 +APERTURE 0.3 OBJECT 0 cube @@ -180,7 +182,6 @@ TRANS -2 5 -2 ROTAT 0 180 0 SCALE 3 3 3 - OBJECT 8 cube material 8 diff --git a/scenes/sphereScene.txt b/scenes/sphereScene.txt new file mode 100644 index 0000000..d7c2dd0 --- /dev/null +++ b/scenes/sphereScene.txt @@ -0,0 +1,128 @@ +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 + +CAMERA +RES 800 800 +FOVY 25 +ITERATIONS 5000 +FILE test.bmp +frame 0 +EYE 2 1 10 +VIEW 0 0 -1 +UP 0 1 0 +FOCAL 9 +APERTURE 0.2 + +OBJECT 0 +cube +material 0 +frame 0 +TRANS 0 0 0 +ROTAT 0 0 0 +SCALE 1 1 1 + diff --git a/src/intersections.h b/src/intersections.h index daefe95..2edbabc 100755 --- a/src/intersections.h +++ b/src/intersections.h @@ -17,8 +17,8 @@ __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__ 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 @@ -70,23 +70,114 @@ __host__ __device__ glm::vec3 getSignOfRay(ray r){ //TODO: IMPLEMENT THIS FUNCTION //Cube intersection test, return -1 if no intersection, otherwise, distance to intersection -__host__ __device__ float boxIntersectionTest(staticGeom box, ray r, glm::vec3& intersectionPoint, glm::vec3& normal){ +__host__ __device__ float boxIntersectionTest(staticGeom box, ray& r, glm::vec3& intersectionPoint, glm::vec3& normal){ - return -1; + //transform ray from world to object space + glm::vec3 ro = multiplyMV(box.inverseTransform, glm::vec4(r.origin, 1.0f)); + glm::vec3 rd = glm::normalize(multiplyMV(box.inverseTransform, glm::vec4(r.direction, 0.0f))); + + ray rt; rt.origin = ro; rt.direction = rd; + + //since cube is in its own unit space, limit between +/- 0.5 + float bL = -0.5f; float bH = 0.5f; + + //check if rd[i] is parallel and outside of any pair of slabs + if(abs(rd.x) <= FLT_EPSILON && ro.x > bH && ro.x < bL) return -1; + if(abs(rd.y) <= FLT_EPSILON && ro.y > bH && ro.y < bL) return -1; + if(abs(rd.z) <= FLT_EPSILON && ro.z > bH && ro.z < bL) return -1; + + //find front and back t-values where ray intersect slab + + //x slabs + float tNear = (bL - ro.x)/rd.x; + float tFar = (bH - ro.x)/rd.x; + + if(tNear > tFar){ + float temp = tFar; + tFar = tNear; + tNear = temp; + } + if(tFar < 0) return -1; + + //y slabs + float tyMin = (bL - ro.y)/rd.y; + float tyMax = (bH - ro.y)/rd.y; + + if(tyMin > tyMax){ + float temp = tyMin; + tyMin = tyMax; + tyMax = temp; + } + if(tyMin > tNear) tNear = tyMin; + if(tyMax < tFar) tFar = tyMax; + + if(tNear > tFar || tFar < 0) return -1; + + //z slabs + float tzMin = (bL - ro.z)/rd.z; + float tzMax = (bH - ro.z)/rd.z; + + if(tzMin > tzMax){ + float temp = tzMin; + tzMin = tzMax; + tzMax = temp; + } + + if(tzMin > tNear) tNear = tzMin; + if(tzMax < tFar) tFar = tzMax; + + if(tNear > tFar || tFar < 0) return -1; + + //if passed all tests, tNear is the intersection + + //find point of intersection and get the normal in object space + glm::vec4 objIntersectionPoint = glm::vec4(getPointOnRay(rt, tNear), 1.0f); + glm::vec4 objNormal(0, 0, 0, 0); //normal in object space + + if(fabs(objIntersectionPoint.x - bH) <= 0.001f){ //if on +x plane, normal is 1 0 0 + objNormal[0] = 1.0f; + } + else if(fabs(objIntersectionPoint.x - bL) <= 0.001f){ //if on -x plane, normal is -1 0 0 + objNormal[0] = -1.0f; + } + else if(fabs(objIntersectionPoint.y - bH) <= 0.001f){ //if on +y plane, normal is 0 1 0 + objNormal[1] = 1.0f; + } + else if(fabs(objIntersectionPoint.y - bL) <= 0.001f){ //if on -y plane, normal is 0 -1 0 + objNormal[1] = -1.0f; + } + else if(fabs(objIntersectionPoint.z - bH) <= 0.001f){ //if on +z plane, normal is 0 0 1 + objNormal[2] = 1.0f; + } + else if(fabs(objIntersectionPoint.z - bL) <= 0.001f){ //if on -z plane, normal is 0 0 -1 + objNormal[2] = -1.0f; + } + + //find intersection point and transform back to world space + glm::vec3 realIntersectionPoint = multiplyMV(box.transform, objIntersectionPoint); + glm::vec3 realOrigin = multiplyMV(box.transform, glm::vec4(0,0,0,1)); + + intersectionPoint = realIntersectionPoint; + normal = glm::normalize(multiplyMV(box.transform, objNormal)); + + return glm::length(intersectionPoint - ro); } //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){ +__host__ __device__ float sphereIntersectionTest(staticGeom sphere, ray& r, glm::vec3& intersectionPoint, glm::vec3& normal){ float radius = .5; - + + //transform ray to object space glm::vec3 ro = multiplyMV(sphere.inverseTransform, glm::vec4(r.origin,1.0f)); glm::vec3 rd = glm::normalize(multiplyMV(sphere.inverseTransform, glm::vec4(r.direction,0.0f))); ray rt; rt.origin = ro; rt.direction = rd; float vDotDirection = glm::dot(rt.origin, rt.direction); + + //discriminant float radicand = vDotDirection * vDotDirection - (glm::dot(rt.origin, rt.origin) - pow(radius, 2)); if (radicand < 0){ return -1; @@ -106,6 +197,7 @@ __host__ __device__ float sphereIntersectionTest(staticGeom sphere, ray r, glm:: t = max(t1, t2); } + //find intersection point and transform back to world space 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)); @@ -177,7 +269,24 @@ __host__ __device__ glm::vec3 getRandomPointOnCube(staticGeom cube, float random //Generates a random point on a given sphere __host__ __device__ glm::vec3 getRandomPointOnSphere(staticGeom sphere, float randomSeed){ - return glm::vec3(0,0,0); + //using trig method from wolfram alpha sphere point picking page + + thrust::default_random_engine rng(hash(randomSeed)); + thrust::uniform_real_distribution u01(-1, 1); + thrust::uniform_real_distribution u02(0, 2*PI); + + glm::vec3 point (0.5f, 0.5f, 0.5f); + + float z = (float)u01(rng); + float theta = (float)u02(rng); + + point.x = sqrt(1 - (z*z)) * cos(theta); + point.y = sqrt( 1 - (z*z)) * sin(theta); + point.z = z; + + glm::vec3 randPoint = multiplyMV(sphere.transform, glm::vec4(point,1.0f)); + + return randPoint; } #endif diff --git a/src/main.cpp b/src/main.cpp index 9b1fdf7..e5b42c5 100755 --- a/src/main.cpp +++ b/src/main.cpp @@ -25,6 +25,10 @@ int main(int argc, char** argv){ bool loadedScene = false; finishedRender = false; + clearImage = false; + + totalTime = 0; + targetFrame = 0; singleFrameMode = false; @@ -66,7 +70,7 @@ int main(int argc, char** argv){ init(argc, argv); #endif - initCuda(); + initCuda(); initVAO(); initTextures(); @@ -123,11 +127,14 @@ void runCuda(){ // execute the kernel - cudaRaytraceCore(dptr, renderCam, targetFrame, iterations, materials, renderScene->materials.size(), geoms, renderScene->objects.size() ); + cudaRaytraceCore(dptr, renderCam, targetFrame, iterations, materials, renderScene->materials.size(), geoms, renderScene->objects.size(), clearImage); // unmap buffer object cudaGLUnmapBufferObject(pbo); }else{ + + //t = clock() - t; + //cout<iterations<iterations; outputImage.setGammaSettings(gamma); string filename = renderCam->imageName; @@ -155,7 +162,7 @@ void runCuda(){ outputImage.saveImageRGB(filename); cout << "Saved frame " << s << " to " << filename << endl; finishedRender = true; - if(singleFrameMode==true){ + if(singleFrameMode==true){ cudaDeviceReset(); exit(0); } @@ -175,6 +182,8 @@ void runCuda(){ } + + #ifdef __APPLE__ void display(){ @@ -199,8 +208,15 @@ void runCuda(){ #else void display(){ + + //t = clock(); + runCuda(); + //t = clock() - t; + //cout<<((float)t)/CLOCKS_PER_SEC<positions[targetFrame].x += 1.0f; + iterations = 0; + clearImage = true; + break; + + case('d'): + renderCam->positions[targetFrame].x -= 1.0f; + iterations = 0; + clearImage = true; + break; + + case('w'): + renderCam->positions[targetFrame].y += 1.0f; + iterations = 0; + clearImage = true; + break; + + case('s'): + renderCam->positions[targetFrame].y -= 1.0f; + iterations = 0; + clearImage = true; + break; + + case('j'): + renderCam->positions[targetFrame].z += 1.0f; + iterations = 0; + clearImage = true; + break; + + case('k'): + renderCam->positions[targetFrame].z -= 1.0f; + iterations = 0; + clearImage = true; + break; + } } diff --git a/src/main.h b/src/main.h index 0bab7cb..51bdcc3 100755 --- a/src/main.h +++ b/src/main.h @@ -30,6 +30,8 @@ #include "utilities.h" #include "scene.h" +#include + #if CUDA_VERSION >= 5000 #include #include @@ -104,6 +106,10 @@ void initTextures(); void initVAO(); GLuint initShader(const char *vertexShaderPath, const char *fragmentShaderPath); +bool clearImage; +clock_t t; +float totalTime; + //------------------------------- //---------CLEANUP STUFF--------- //------------------------------- diff --git a/src/raytraceKernel.cu b/src/raytraceKernel.cu index b4f4ec5..3e44293 100755 --- a/src/raytraceKernel.cu +++ b/src/raytraceKernel.cu @@ -44,10 +44,36 @@ __host__ __device__ glm::vec3 generateRandomNumberFromThread(glm::vec2 resolutio //TODO: IMPLEMENT THIS FUNCTION //Function that does the initial raycast from the camera __host__ __device__ ray raycastFromCameraKernel(glm::vec2 resolution, float time, int x, int y, glm::vec3 eye, glm::vec3 view, glm::vec3 up, glm::vec2 fov){ - ray r; - r.origin = glm::vec3(0,0,0); - r.direction = glm::vec3(0,0,-1); - return r; + + ray r; + + //ray creation from camear stuff to be used in raycastFromCameraKernel + glm::vec3 M = eye + view; //center of screen + + //project screen to world space + glm::vec3 A = glm::cross(view, up); + glm::vec3 B = glm::cross(A, view); + + float C = glm::length(view); + + float phi = fov.y/180.0f * PI; //convert to radians + B = glm::normalize(B); + glm::vec3 V = C * tan(phi) * B; + + float theta = fov.x/180.0f * PI; + A = glm::normalize(A); + glm::vec3 H = C * tan(theta) * A; + + //find the world space coord of the pixel + float sx = (float)x / (resolution.x-1.0f); + float sy = (float)y / (resolution.y-1.0f); + + glm::vec3 P = M + H * (2.0f * sx - 1.0f) + V * (1.0f - 2.0f * sy); + + r.origin = eye; + r.direction = glm::normalize(P - eye); + + return r; } //Kernel that blacks out a given image buffer @@ -61,18 +87,20 @@ __global__ void clearImage(glm::vec2 resolution, glm::vec3* image){ } //Kernel that writes the image to the OpenGL PBO directly. -__global__ void sendImageToPBO(uchar4* PBOpos, glm::vec2 resolution, glm::vec3* image){ +__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); + //iterations = 1; + 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; + color.x = image[index].x*255.0 / iterations; + color.y = image[index].y*255.0 / iterations; + color.z = image[index].z*255.0 / iterations; if(color.x>255){ color.x = 255; @@ -94,30 +122,195 @@ __global__ void sendImageToPBO(uchar4* PBOpos, glm::vec2 resolution, glm::vec3* } } + +//Does intersection on all of the objects and returns length of closest intersection +__host__ __device__ float testGeomIntersection(staticGeom* geoms, int numberOfGeoms, ray& r, glm::vec3& intersectionPoint, glm::vec3& normal, int& objID){ + + float len = FLT_MAX; + float tempLen = -1; + glm::vec3 tempIntersection; + glm::vec3 tempNormal; + + //check for interesction + for(int geomInd = 0; geomInd 0){ + colors[index] += surfColor; + return; + } + +#pragma region lightAndShadow + + glm::vec3 diffuse(0,0,0); + glm::vec3 phong(0,0,0); + + //do light and shadow computation + for(int i = 0; i < numLights; ++i){ + + int lightGeomID = lightID[i]; + glm::vec3 lightPos; + glm::vec3 lightColor = materials[geoms[lightGeomID].materialid].color; + + //find a random point on the light + if(geoms[lightGeomID].type == CUBE){ + lightPos = getRandomPointOnCube(geoms[lightGeomID], time); //CHANGE TO TIME! + } + else if(geoms[lightGeomID].type == SPHERE){ + lightPos = getRandomPointOnSphere(geoms[lightGeomID], time); //CHANGE TO TIME! + } + + //find vector from intersection to point on light + glm::vec3 L = lightPos - intersection; + float distToLight = glm::length(L); + L = glm::normalize(L); + + //check if in shadow + objID = -1; + ray shadowFeeler; + shadowFeeler.direction = L; + shadowFeeler.origin = intersection + 0.0001f*L; //offset origin a little bit so it doesn't self intersect + + glm::vec3 shadowNormal; glm::vec3 shadowIntersection; + len = testGeomIntersection(geoms, numberOfGeoms, shadowFeeler, shadowIntersection, shadowNormal, objID); + + //if intersection occured and intersection is in between the intersection point and the light position + if(objID != -1 && len < distToLight){ + + if(materials[geoms[objID].materialid].emittance == 0){ //only cast shadow if we intersected with object that is not a light + //color is ambient color + finalColor = glm::vec3(0,0,0); + continue; + } + } + + //do diffuse calculation + diffuse += glm::clamp(glm::dot(L, normal), 0.0f, 1.0f) * surfColor * lightColor; + + //clamp diffuse to surface color + diffuse.x = clamp(diffuse.x, 0.0f, surfColor.x); + diffuse.y = clamp(diffuse.y, 0.0f, surfColor.y); + diffuse.z = clamp(diffuse.z, 0.0f, surfColor.z); + + //specular + if(materials[matID].specularExponent != 0){ + glm::vec3 R = glm::normalize( -L - 2.0f*glm::dot(-L, normal) *normal); + glm::vec3 V = -firstRay.direction; //already normalized + + phong += materials[matID].specularColor * + pow(glm::clamp(glm::dot(R, V), 0.0f, 1.0f), materials[matID].specularExponent) * lightColor; + + //phong *= 0.5f; + //diffuse *= 0.9f; + } + } + +#pragma endregion lightAndShadow + + finalColor += glm::clamp(diffuse + phong, 0.0f, 1.0f); + //output final color + colors[index] += finalColor; + + } } + //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){ +void cudaRaytraceCore(uchar4* PBOpos, camera* renderCam, int frame, int iterations, material* materials, int numberOfMaterials, geom* geoms, int numberOfGeoms, bool& clear){ int traceDepth = 1; //determines how many bounces the raytracer traces // set up crucial magic int tileSize = 8; - dim3 threadsPerBlock(tileSize, tileSize); + dim3 threadsPerBlock(tileSize, tileSize); //each block has 8 * 8 threads dim3 fullBlocksPerGrid((int)ceil(float(renderCam->resolution.x)/float(tileSize)), (int)ceil(float(renderCam->resolution.y)/float(tileSize))); //send image to GPU @@ -125,6 +318,8 @@ void cudaRaytraceCore(uchar4* PBOpos, camera* renderCam, int frame, int iteratio cudaMalloc((void**)&cudaimage, (int)renderCam->resolution.x*(int)renderCam->resolution.y*sizeof(glm::vec3)); cudaMemcpy( cudaimage, renderCam->image, (int)renderCam->resolution.x*(int)renderCam->resolution.y*sizeof(glm::vec3), cudaMemcpyHostToDevice); + std::vector lightVec; + //package geometry and materials and sent to GPU staticGeom* geomList = new staticGeom[numberOfGeoms]; for(int i=0; i 0) + lightVec.push_back(i); + } staticGeom* cudageoms = NULL; cudaMalloc((void**)&cudageoms, numberOfGeoms*sizeof(staticGeom)); cudaMemcpy( cudageoms, geomList, numberOfGeoms*sizeof(staticGeom), cudaMemcpyHostToDevice); + + //copy materials to memory + material* cudaMaterials = NULL; + cudaMalloc((void**)&cudaMaterials, numberOfMaterials*sizeof(material)); + cudaMemcpy( cudaMaterials, materials, numberOfMaterials*sizeof(material), cudaMemcpyHostToDevice); + + //copy light ID to memeory + int numLights = lightVec.size(); + int* lightID = new int[numLights]; + for(int i = 0; i views[frame]; cam.up = renderCam->ups[frame]; cam.fov = renderCam->fov; + cam.focalLength = renderCam->focalLengths[frame]; + cam.aperture = renderCam->apertures[frame]; //kernel launches - raytraceRay<<>>(renderCam->resolution, (float)iterations, cam, traceDepth, cudaimage, cudageoms, numberOfGeoms); + if(clear){ + clearImage<<>>(renderCam->resolution, cudaimage); + clear = false; + } + else + raytraceRay<<>>(renderCam->resolution, (float)iterations, cam, traceDepth, cudaimage, cudageoms, numberOfGeoms, + cudaMaterials, numLights, cudaLights); - sendImageToPBO<<>>(PBOpos, renderCam->resolution, cudaimage); + 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); @@ -162,10 +385,13 @@ void cudaRaytraceCore(uchar4* PBOpos, camera* renderCam, int frame, int iteratio //free up stuff, or else we'll leak memory like a madman cudaFree( cudaimage ); cudaFree( cudageoms ); + cudaFree( cudaMaterials); delete geomList; + delete lightID; // make certain the kernel has completed cudaThreadSynchronize(); checkCUDAError("Kernel failed!"); } + diff --git a/src/raytraceKernel.h b/src/raytraceKernel.h index 5fcf5a3..41e4339 100755 --- a/src/raytraceKernel.h +++ b/src/raytraceKernel.h @@ -20,6 +20,7 @@ #include #endif -void cudaRaytraceCore(uchar4* pos, camera* renderCam, int frame, int iterations, material* materials, int numberOfMaterials, geom* geoms, int numberOfGeoms); +void cudaRaytraceCore(uchar4* pos, camera* renderCam, int frame, int iterations, material* materials, int numberOfMaterials, geom* geoms, int numberOfGeoms, bool& clear); + #endif diff --git a/src/scene.cpp b/src/scene.cpp index 415d627..95090c5 100755 --- a/src/scene.cpp +++ b/src/scene.cpp @@ -162,6 +162,8 @@ int scene::loadCamera(){ vector positions; vector views; vector ups; + vector focalLengths; + vector apertures; while (!line.empty() && fp_in.good()){ //check frame number @@ -172,7 +174,7 @@ int scene::loadCamera(){ } //load camera properties - for(int i=0; i<3; i++){ + for(int i=0; i<5; i++){ //glm::vec3 translation; glm::vec3 rotation; glm::vec3 scale; utilityCore::safeGetline(fp_in,line); tokens = utilityCore::tokenizeString(line); @@ -183,6 +185,13 @@ int scene::loadCamera(){ }else if(strcmp(tokens[0].c_str(), "UP")==0){ ups.push_back(glm::vec3(atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str()))); } + else if(strcmp(tokens[0].c_str(), "FOCAL") ==0){ + focalLengths.push_back(atof(tokens[1].c_str())); + } + else if(strcmp(tokens[0].c_str(), "APERTURE") ==0){ + apertures.push_back(atof(tokens[1].c_str())); + } + } frameCount++; @@ -194,10 +203,15 @@ int scene::loadCamera(){ newCamera.positions = new glm::vec3[frameCount]; newCamera.views = new glm::vec3[frameCount]; newCamera.ups = new glm::vec3[frameCount]; + newCamera.focalLengths = new float[frameCount]; + newCamera.apertures = new float[frameCount]; + for(int i=0; i