diff --git a/ImplementationDetails.txt b/ImplementationDetails.txt new file mode 100644 index 0000000..011e3a4 --- /dev/null +++ b/ImplementationDetails.txt @@ -0,0 +1,15 @@ +Implemented the following: + +* Depth of field with aperture control (Make the aperture value in the scene files greater than 25 to have pin hole effect) +* Mesh (obj) loader +* Full reflection and refraction using Fresnel equations +* Tried some hard coded motion blur, but didn't get time do it using the frame values +* Coded up the Cook Torrance BRDF, but did not get time to tweak it + +Code: +https://github.com/tijutv/Project2-Pathtracer + +Blog: +http://cudapathtracer.blogspot.com/ + + diff --git a/PROJ1_WIN/565Raytracer.suo b/PROJ1_WIN/565Raytracer.suo new file mode 100644 index 0000000..9a2d185 Binary files /dev/null and b/PROJ1_WIN/565Raytracer.suo differ diff --git a/PROJ1_WIN/565Raytracer/565Raytracer.vcxproj b/PROJ1_WIN/565Raytracer/565Raytracer.vcxproj index fcc853d..5e86015 100755 --- a/PROJ1_WIN/565Raytracer/565Raytracer.vcxproj +++ b/PROJ1_WIN/565Raytracer/565Raytracer.vcxproj @@ -1,126 +1,128 @@ - - - - - Debug - Win32 - - - Release - Win32 - - - - - - - - - - - - - - - - - - - - - - - - - - - - - - {FF21CA49-522E-4E86-B508-EE515B248FC4} - Win32Proj - 565Raytracer - 565Raytracer - - - - Application - true - Unicode - - - Application - false - true - Unicode - - - - - - - - - - - - - - true - - - false - - - - - - Level3 - Disabled - WIN32;_DEBUG;_CONSOLE;%(PreprocessorDefinitions) - C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.0\include;C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\C\common\inc;../shared/glew/include;../shared/freeglut/include;%(AdditionalIncludeDirectories) - - - Console - true - ../shared/glew/lib;../shared/freeglut/lib;%(AdditionalLibraryDirectories) - cudart.lib; glew32.lib;glu32.lib;opengl32.lib;kernel32.lib;user32.lib;gdi32.lib;winspool.lib;comdlg32.lib;advapi32.lib;shell32.lib;ole32.lib;oleaut32.lib;uuid.lib;odbc32.lib;odbccp32.lib;%(AdditionalDependencies) - mainCRTStartup - - - - - $(ProjectDir)$(Platform)/$(Configuration)/%(Filename)%(Extension).obj - C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.0\include;C:/ProgramData/NVIDIA Corporation/NVIDIA GPU Computing SDK 4.0/C/common/inc;../shared/glew/includes;../shared/freeglut/includes - - - - - Level3 - - - MaxSpeed - true - true - WIN32;NDEBUG;_CONSOLE;%(PreprocessorDefinitions) - C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.0\include;C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\C\common\inc;../shared/glew/include;../shared/freeglut/include;%(AdditionalIncludeDirectories) - - - Console - true - true - true - ../shared/glew/lib;../shared/freeglut/lib;%(AdditionalLibraryDirectories) - cudart.lib; glew32.lib;glu32.lib;opengl32.lib;kernel32.lib;user32.lib;gdi32.lib;winspool.lib;comdlg32.lib;advapi32.lib;shell32.lib;ole32.lib;oleaut32.lib;uuid.lib;odbc32.lib;odbccp32.lib;%(AdditionalDependencies) - mainCRTStartup - - - $(ProjectDir)$(Platform)/$(Configuration)/%(Filename)%(Extension).obj - C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.0\include;C:/ProgramData/NVIDIA Corporation/NVIDIA GPU Computing SDK 4.0/C/common/inc;../shared/glew/includes;../shared/freeglut/includes - - - - - - + + + + + Debug + Win32 + + + Release + Win32 + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + {FF21CA49-522E-4E86-B508-EE515B248FC4} + Win32Proj + 565Raytracer + 565Raytracer + + + + Application + true + Unicode + + + Application + false + true + Unicode + + + + + + + + + + + + + + true + + + false + + + + + + Level3 + Disabled + WIN32;_DEBUG;_CONSOLE;%(PreprocessorDefinitions) + C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.0\include;C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\C\common\inc;../shared/glew/include;../shared/freeglut/include;%(AdditionalIncludeDirectories) + + + Console + true + ../shared/glew/lib;../shared/freeglut/lib;%(AdditionalLibraryDirectories) + cudart.lib; glew32.lib;glu32.lib;opengl32.lib;kernel32.lib;user32.lib;gdi32.lib;winspool.lib;comdlg32.lib;advapi32.lib;shell32.lib;ole32.lib;oleaut32.lib;uuid.lib;odbc32.lib;odbccp32.lib;%(AdditionalDependencies) + mainCRTStartup + + + + + $(ProjectDir)$(Platform)/$(Configuration)/%(Filename)%(Extension).obj + C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.0\include;C:/ProgramData/NVIDIA Corporation/NVIDIA GPU Computing SDK 4.0/C/common/inc;../shared/glew/includes;../shared/freeglut/includes + compute_20,sm_20 + + + + + Level3 + + + MaxSpeed + true + true + WIN32;NDEBUG;_CONSOLE;%(PreprocessorDefinitions) + C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.0\include;C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\C\common\inc;../shared/glew/include;../shared/freeglut/include;%(AdditionalIncludeDirectories) + + + Console + true + true + true + ../shared/glew/lib;../shared/freeglut/lib;%(AdditionalLibraryDirectories) + cudart.lib; glew32.lib;glu32.lib;opengl32.lib;kernel32.lib;user32.lib;gdi32.lib;winspool.lib;comdlg32.lib;advapi32.lib;shell32.lib;ole32.lib;oleaut32.lib;uuid.lib;odbc32.lib;odbccp32.lib;%(AdditionalDependencies) + mainCRTStartup + + + $(ProjectDir)$(Platform)/$(Configuration)/%(Filename)%(Extension).obj + C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.0\include;C:/ProgramData/NVIDIA Corporation/NVIDIA GPU Computing SDK 4.0/C/common/inc;../shared/glew/includes;../shared/freeglut/includes + compute_20,sm_20 + + + + + + \ No newline at end of file diff --git a/PROJ1_WIN/565Raytracer/565Raytracer.vcxproj.user b/PROJ1_WIN/565Raytracer/565Raytracer.vcxproj.user index d7ca222..9334dc8 100755 --- a/PROJ1_WIN/565Raytracer/565Raytracer.vcxproj.user +++ b/PROJ1_WIN/565Raytracer/565Raytracer.vcxproj.user @@ -1,7 +1,11 @@ - - - - scene="../../scenes/sampleScene.txt" - WindowsLocalDebugger - + + + + scene="../../scenes/sampleScene_1.txt" + WindowsLocalDebugger + + + scene="../../scenes/sampleScene_2.txt" + WindowsLocalDebugger + \ No newline at end of file diff --git a/PROJ1_WIN/565Raytracer/Mesh/tetrahedron.obj b/PROJ1_WIN/565Raytracer/Mesh/tetrahedron.obj new file mode 100644 index 0000000..c3d6c09 --- /dev/null +++ b/PROJ1_WIN/565Raytracer/Mesh/tetrahedron.obj @@ -0,0 +1,61 @@ +# OBJ file created by ply_to_obj.c +# +g Object001 + +v -0.57735 -0.57735 0.57735 +v 0.934172 0.356822 0 +v 0.934172 -0.356822 0 +v -0.934172 0.356822 0 +v -0.934172 -0.356822 0 +v 0 0.934172 0.356822 +v 0 0.934172 -0.356822 +v 0.356822 0 -0.934172 +v -0.356822 0 -0.934172 +v 0 -0.934172 -0.356822 +v 0 -0.934172 0.356822 +v 0.356822 0 0.934172 +v -0.356822 0 0.934172 +v 0.57735 0.57735 -0.57735 +v 0.57735 0.57735 0.57735 +v -0.57735 0.57735 -0.57735 +v -0.57735 0.57735 0.57735 +v 0.57735 -0.57735 -0.57735 +v 0.57735 -0.57735 0.57735 +v -0.57735 -0.57735 -0.57735 + +f 19 3 2 +f 12 19 2 +f 15 12 2 +f 8 14 2 +f 18 8 2 +f 3 18 2 +f 20 5 4 +f 9 20 4 +f 16 9 4 +f 13 17 4 +f 1 13 4 +f 5 1 4 +f 7 16 4 +f 6 7 4 +f 17 6 4 +f 6 15 2 +f 7 6 2 +f 14 7 2 +f 10 18 3 +f 11 10 3 +f 19 11 3 +f 11 1 5 +f 10 11 5 +f 20 10 5 +f 20 9 8 +f 10 20 8 +f 18 10 8 +f 9 16 7 +f 8 9 7 +f 14 8 7 +f 12 15 6 +f 13 12 6 +f 17 13 6 +f 13 1 11 +f 12 13 11 +f 19 12 11 diff --git a/PROJ1_WIN/565Raytracer/renders/DOF.PNG b/PROJ1_WIN/565Raytracer/renders/DOF.PNG new file mode 100644 index 0000000..a55a2d4 Binary files /dev/null and b/PROJ1_WIN/565Raytracer/renders/DOF.PNG differ diff --git a/PROJ1_WIN/565Raytracer/renders/DOF1.PNG b/PROJ1_WIN/565Raytracer/renders/DOF1.PNG new file mode 100644 index 0000000..04763bd Binary files /dev/null and b/PROJ1_WIN/565Raytracer/renders/DOF1.PNG differ diff --git a/PROJ1_WIN/565Raytracer/renders/DOF_18.bmp b/PROJ1_WIN/565Raytracer/renders/DOF_18.bmp new file mode 100644 index 0000000..ab359bd Binary files /dev/null and b/PROJ1_WIN/565Raytracer/renders/DOF_18.bmp differ diff --git a/PROJ1_WIN/565Raytracer/renders/DOF_1_8.bmp b/PROJ1_WIN/565Raytracer/renders/DOF_1_8.bmp new file mode 100644 index 0000000..4a684a6 Binary files /dev/null and b/PROJ1_WIN/565Raytracer/renders/DOF_1_8.bmp differ diff --git a/PROJ1_WIN/565Raytracer/renders/DOF_8.bmp b/PROJ1_WIN/565Raytracer/renders/DOF_8.bmp new file mode 100644 index 0000000..95a05d3 Binary files /dev/null and b/PROJ1_WIN/565Raytracer/renders/DOF_8.bmp differ diff --git a/PROJ1_WIN/565Raytracer/renders/DOF_Infinity.bmp b/PROJ1_WIN/565Raytracer/renders/DOF_Infinity.bmp new file mode 100644 index 0000000..53e8bac Binary files /dev/null and b/PROJ1_WIN/565Raytracer/renders/DOF_Infinity.bmp differ diff --git a/PROJ1_WIN/565Raytracer/renders/FirstBlogImage.bmp b/PROJ1_WIN/565Raytracer/renders/FirstBlogImage.bmp new file mode 100644 index 0000000..679ce7d Binary files /dev/null and b/PROJ1_WIN/565Raytracer/renders/FirstBlogImage.bmp differ diff --git a/PROJ1_WIN/565Raytracer/renders/Mesh_MotionBlur.bmp b/PROJ1_WIN/565Raytracer/renders/Mesh_MotionBlur.bmp new file mode 100644 index 0000000..46f28de Binary files /dev/null and b/PROJ1_WIN/565Raytracer/renders/Mesh_MotionBlur.bmp differ diff --git a/PROJ1_WIN/565Raytracer/renders/sampleScene.0.bmp b/PROJ1_WIN/565Raytracer/renders/sampleScene.0.bmp new file mode 100644 index 0000000..4b6de2f Binary files /dev/null and b/PROJ1_WIN/565Raytracer/renders/sampleScene.0.bmp differ diff --git a/PROJ1_WIN/565Raytracer/renders/sampleScene.0_temp.bmp b/PROJ1_WIN/565Raytracer/renders/sampleScene.0_temp.bmp new file mode 100644 index 0000000..9828781 Binary files /dev/null and b/PROJ1_WIN/565Raytracer/renders/sampleScene.0_temp.bmp differ diff --git a/README.md b/README.md index cbc1dce..c43e9a4 100755 --- a/README.md +++ b/README.md @@ -1,10 +1,32 @@ + ------------------------------------------------------------------------------- CIS565: Project 2: CUDA Pathtracer ------------------------------------------------------------------------------- Fall 2012 ------------------------------------------------------------------------------- -Due Friday, 10/12/2012 + +------------------------------------------------------------------------------- +IMPLEMENTATION: ------------------------------------------------------------------------------- +* Depth of field with aperture control (Make the aperture value in the scene files greater than 25 to have pin hole effect) +* Mesh (obj) loader +* Full reflection and refraction using Fresnel equations +* Tried some hard coded motion blur, but didn't get time do it using the frame values +* Coded up the Cook Torrance BRDF, but did not get time to tweak it + +------------------------------------------------------------------------------- +CODE: +------------------------------------------------------------------------------- +https://github.com/tijutv/Project2-Pathtracer + +------------------------------------------------------------------------------- +BLOG: +------------------------------------------------------------------------------- +http://cudapathtracer.blogspot.com/ + +------------------------------------------------------------------------------- + + ------------------------------------------------------------------------------- NOTE: diff --git a/ScreenCapture.PNG b/ScreenCapture.PNG new file mode 100644 index 0000000..98d8c83 Binary files /dev/null and b/ScreenCapture.PNG differ diff --git a/ScreenCapture_10-12-2012 10.41.23 PM.wmv b/ScreenCapture_10-12-2012 10.41.23 PM.wmv new file mode 100644 index 0000000..562a549 Binary files /dev/null and b/ScreenCapture_10-12-2012 10.41.23 PM.wmv differ diff --git a/scenes/sampleScene_1.txt b/scenes/sampleScene_1.txt new file mode 100644 index 0000000..4b0800e --- /dev/null +++ b/scenes/sampleScene_1.txt @@ -0,0 +1,292 @@ +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 .6 .6 .04 +SPECEX 1 +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 1.5 +SCATTER 0 +ABSCOEFF .02 5.1 5.7 +RSCTCOEFF 13 +EMITTANCE 0 + +MATERIAL 6 //green glossy +RGB .3 .3 .3 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0.2 +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 9 + +MATERIAL 9 //mirror +RGB 0 0 0 +SPECEX 0 +SPECRGB 1 1 1 +REFL 1 +REFR 0 +REFRIOR 1.5 +SCATTER 0 +ABSCOEFF .02 5.1 5.7 +RSCTCOEFF 13 +EMITTANCE 0 + +MATERIAL 10 //outside light +RGB 0.2 0.2 0.2 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 1 + +MATERIAL 11 //blue diffuse +RGB .03 .46 .84 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 0 +REFRIOR 0 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 0 + +MATERIAL 12 //blue diffuse +RGB .43 .46 .14 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 0 +REFRIOR 0 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 0 + +MATERIAL 13 //blue diffuse +RGB .43 .16 .54 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 0 +REFRIOR 0 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 0 + +CAMERA +RES 800 800 +FOVY 18 +ITERATIONS 3000 +FILE renders/sampleScene.bmp +APERTURE 2.8 +FOCAL_DISTANCE 20 +frame 0 +EYE 0 5 20 +VIEW 0 0 -1 +UP 0 1 0 +frame 1 +EYE 0 5 20 +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 +frame 1 +TRANS 0 0 0 +ROTAT 0 0 90 +SCALE .01 10 10 + +OBJECT 1 +cube +material 11 +frame 0 +TRANS 0 5 -5 +ROTAT 0 90 0 +SCALE .01 10 10 +frame 1 +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 +frame 1 +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 +frame 1 +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 +frame 1 +TRANS 5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +OBJECT 5 +sphere +material 0 +frame 0 +TRANS 3 3.5 0 +ROTAT 0 180 0 +SCALE 2 2 2 +frame 1 +TRANS 3 1.5 3 +ROTAT 0 180 0 +SCALE 2 2 2 + + +OBJECT 6 +sphere +material 12 +frame 0 +TRANS -1 2 -1 +ROTAT 0 180 0 +SCALE 3 3 3 +frame 1 +TRANS -1 2 0 +ROTAT 0 180 0 +SCALE 3 3 3 + +OBJECT 7 +cube +material 8 +frame 0 +TRANS 0 10 0 +ROTAT 0 0 90 +SCALE .3 5 5 +frame 1 +TRANS 0 10 0 +ROTAT 0 0 90 +SCALE .3 5 5 + +OBJECT 8 +Mesh\tetrahedron.obj +material 13 +frame 0 +TRANS -2 6 0 +ROTAT 0 0 90 +SCALE 1.5 1.5 1.5 +frame 1 +TRANS -2 5 0 +ROTAT 0 0 90 +SCALE 1.5 1.5 1.5 diff --git a/scenes/tetrahedron.obj b/scenes/tetrahedron.obj new file mode 100644 index 0000000..c3d6c09 --- /dev/null +++ b/scenes/tetrahedron.obj @@ -0,0 +1,61 @@ +# OBJ file created by ply_to_obj.c +# +g Object001 + +v -0.57735 -0.57735 0.57735 +v 0.934172 0.356822 0 +v 0.934172 -0.356822 0 +v -0.934172 0.356822 0 +v -0.934172 -0.356822 0 +v 0 0.934172 0.356822 +v 0 0.934172 -0.356822 +v 0.356822 0 -0.934172 +v -0.356822 0 -0.934172 +v 0 -0.934172 -0.356822 +v 0 -0.934172 0.356822 +v 0.356822 0 0.934172 +v -0.356822 0 0.934172 +v 0.57735 0.57735 -0.57735 +v 0.57735 0.57735 0.57735 +v -0.57735 0.57735 -0.57735 +v -0.57735 0.57735 0.57735 +v 0.57735 -0.57735 -0.57735 +v 0.57735 -0.57735 0.57735 +v -0.57735 -0.57735 -0.57735 + +f 19 3 2 +f 12 19 2 +f 15 12 2 +f 8 14 2 +f 18 8 2 +f 3 18 2 +f 20 5 4 +f 9 20 4 +f 16 9 4 +f 13 17 4 +f 1 13 4 +f 5 1 4 +f 7 16 4 +f 6 7 4 +f 17 6 4 +f 6 15 2 +f 7 6 2 +f 14 7 2 +f 10 18 3 +f 11 10 3 +f 19 11 3 +f 11 1 5 +f 10 11 5 +f 20 10 5 +f 20 9 8 +f 10 20 8 +f 18 10 8 +f 9 16 7 +f 8 9 7 +f 14 8 7 +f 12 15 6 +f 13 12 6 +f 17 13 6 +f 13 1 11 +f 12 13 11 +f 19 12 11 diff --git a/src/interactions.h b/src/interactions.h index e18cfff..7dea481 100755 --- a/src/interactions.h +++ b/src/interactions.h @@ -1,106 +1,340 @@ -// 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; +}; + +//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, bool& internalReflection); +__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, float xi3); + +void GetParametersForRayCast(glm::vec2 resolution, glm::vec3 eye, glm::vec3 view, glm::vec3 up, + glm::vec2 fov, glm::vec3& M, glm::vec3& A, glm::vec3& B, float& distImagePlaneFromCamera){ + ray r; + r.origin = eye; + view = glm::normalize(view); + up = glm::normalize(up); + + A = glm::normalize(glm::cross(view, up)); + B = glm::normalize(glm::cross(A, view)); + + float tanVert = tan(fov.y*PI/180); + float tanHor = tan(fov.x*PI/180); + + float camDistFromScreen = (float)((resolution.y/2.0)/tanVert); + glm::vec3 C = view*camDistFromScreen; + M = eye + C; + + distImagePlaneFromCamera = camDistFromScreen; +} + +//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, bool& internalReflection) { + float n12 = incidentIOR/transmittedIOR; + + float cos1 = glm::dot(normal, glm::vec3(-incident.x, -incident.y, -incident.z)); + float rootValue = 1 - (n12*n12)*(1.0f-cos1*cos1); + if (rootValue < 0) + { + // Internal Reflection + internalReflection = true; + return calculateReflectionDirection(normal, incident); + } + + if (cos1 > 0.0) + return glm::normalize(normal*(n12*cos1 - sqrt(rootValue)) + incident*n12); + else + return glm::normalize(normal*(-n12*cos1 + sqrt(rootValue)) + incident*n12); +} + +//TODO (OPTIONAL): IMPLEMENT THIS FUNCTION +__host__ __device__ glm::vec3 calculateReflectionDirection(glm::vec3 normal, glm::vec3 incident) { + //nothing fancy here + // Rr = Ri - 2N(Ri.N) + float dotProd = glm::dot(incident, normal); + return glm::normalize(incident - normal*2.0f*dotProd); +} + +//TODO (OPTIONAL): IMPLEMENT THIS FUNCTION +__host__ __device__ Fresnel calculateFresnel(glm::vec3 normal, glm::vec3 incident, float incidentIOR, float transmittedIOR) { + Fresnel fresnel; + + incident = glm::normalize(incident); + normal = glm::normalize(normal); + float cosIncidence = abs(glm::dot(incident, normal)); + float sinIncidence = sqrt(1-cosIncidence*cosIncidence); + + if (transmittedIOR > 0.0 && incidentIOR > 0) + { + float sinRefract = (incidentIOR/transmittedIOR)*sinIncidence; + float commonNumerator = sqrt(1-sinRefract*sinRefract); + float RsNumerator = incidentIOR*cosIncidence-transmittedIOR*commonNumerator; + float RsDenominator = incidentIOR*cosIncidence+transmittedIOR*commonNumerator; + float Rs = 0; + if (RsDenominator != 0) + Rs = (RsNumerator/RsDenominator)*(RsNumerator/RsDenominator); + + float RpNumerator = (incidentIOR * commonNumerator) - (transmittedIOR * cosIncidence); + float RpDenominator = (incidentIOR * commonNumerator) + (transmittedIOR * cosIncidence); + float Rp = 0; + if (RpDenominator != 0) + Rp = (RpNumerator/RpDenominator)*(RpNumerator/RpDenominator); + + fresnel.reflectionCoefficient = (Rs + Rp)/2.0; + fresnel.transmissionCoefficient = 1 - fresnel.reflectionCoefficient; + } + else + { + 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, float xi3) { + + //crucial difference between this and calculateRandomDirectionInSphere: THIS IS COSINE WEIGHTED! + + float up = sqrt(xi1); // cos(theta) + float over = sqrt(1 - up * up); // sin(theta) + if (xi3 <= 0.5) + over = -over; + 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 = glm::vec3(0,0,0); + if (abs(normal.x) < SQRT_OF_ONE_THIRD) { + directionNotNormal.x = 1; + } else if (abs(normal.y) < SQRT_OF_ONE_THIRD) { + directionNotNormal.y = 1; + } else { + directionNotNormal.z = 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)); + + glm::vec3 direction = glm::vec3(( up * normal ) + ( sin(around) * over * perpendicularDirection1 ) + ( cos(around) * over * perpendicularDirection2 )); + direction = glm::normalize(direction); + if (glm::dot(direction,normal)<0) + direction = -direction; + return direction; + + /*float theta = acos(sqrt(1.0-xi1)); + float phi = TWO_PI * xi2; + + float xs = sinf(theta) * cosf(phi); + float ys = cosf(theta); + float zs = sinf(theta) * sinf(phi); + + glm::vec3 y = glm::vec3(normal.x, normal.y, normal.z); + glm::vec3 h = y; + if (fabs(h.x)<=fabs(h.y) && fabs(h.x)<=fabs(h.z)) + h.x= 1.0; + else if (fabs(h.y)<=fabs(h.x) && fabs(h.y)<=fabs(h.z)) + h.y= 1.0; + else + h.z= 1.0; + + + glm::vec3 x = glm::normalize(glm::cross(h, y)); + glm::vec3 z = glm::normalize(glm::cross(x, y)); + + glm::vec3 direction = xs * x + ys * y + zs * z; + direction = glm::normalize(direction); + if (glm::dot(direction,normal)<0) + direction = -direction; + return direction;*/ + +} + +//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) { + float z = 1.0f - 2.0f*xi1; + float temp = 1.0f - z*z; + float r; + if (temp < 0.0f) + r = 0.0f; + else + r = sqrtf(temp); + + float phi = 2.0f * PI * xi2; + float x = r * cosf(phi); + float y = r * sinf(phi); + return glm::normalize(glm::vec3(x, y, z)); +} + +//__host__ __device__ void CookTorrence(ray inRay, ray& outRay, glm::vec3 intersect, glm::vec3 N, glm::vec3 emittedColor, +// glm::vec3& color, material mat, bool inside, float xi1, float xi2) +//{ +// // Reference: http://renderman.pixar.com/view/cook-torrance-shader +// float gaussConstant = 100; +// float m = 0.2; // roughness +// +// // Calculate reflective and transmittive coefficients using Fresnel equations +// float n1, n2; +// if (inside) +// { +// n2 = 1.0f; +// n1 = mat.indexOfRefraction; +// } +// else +// { +// n1 = 1.0f; +// n2 = mat.indexOfRefraction; +// } +// +// Fresnel fresnel = calculateFresnel(N, intersect, n1, n2); +// outRay.direction = calculateRandomDirectionInHemisphere(N, xi1, xi2); +// glm::vec3 L = outRay.direction; +// +// glm::vec3 V = glm::normalize(inRay.direction * -1.0f); +// +// // Half angle vector +// glm::vec3 H = glm::normalize(V+L); +// +// // Attenuation Factor G +// float NDotV = glm::dot(N, V); +// float NDotH = glm::dot(N, H); +// float VDotH = glm::dot(V, H); +// float NDotL = glm::dot(N, L); +// +// float G = glm::min(2*NDotH*NDotV/VDotH, 2*NDotH*NDotL/VDotH); +// G = glm::min(1.0f, G); +// +// // Microfacet Slope Distribution D +// float alpha = acos(NDotH); +// float D = gaussConstant*glm::exp(-(alpha*alpha)/(m*m)); +// float cook = (fresnel.reflectionCoefficient*D*G)/(PI*NDotL*NDotV); +// color = emittedColor * mat.color;// * cook; +//} + +//TODO (PARTIALLY OPTIONAL): IMPLEMENT THIS FUNCTION +//returns 0 if diffuse scatter, 1 if reflected, 2 if transmitted. +__host__ __device__ int calculateBSDF(ray inRay, ray& outRay, glm::vec3 intersect, glm::vec3 normal, glm::vec3 emittedColor, + /*AbsorptionAndScatteringProperties& currentAbsorptionAndScattering, */ + glm::vec3& color, /*glm::vec3& unabsorbedColor,*/ material m, bool inside, float xi1, float xi2, float xi3){ + outRay.origin = intersect; + + if (m.reflectivity >= (1.0f-0.01f)) + { + // Specular reflection + outRay.direction = calculateReflectionDirection(normal, inRay.direction); + color = emittedColor; + return 1; + } + else if (m.hasRefractive) + { + // Calculate reflective and transmittive coefficients using Fresnel equations + float n1, n2; + if (inside) + { + n2 = 1.0f; + n1 = m.indexOfRefraction; + + // Do refraction + bool internalReflection = false; + outRay.direction = calculateTransmissionDirection(normal, inRay.direction, n1, n2, internalReflection); + color = emittedColor;//*fresnel.transmissionCoefficient; + if (internalReflection) + return 1; + else + return 2; + } + else + { + n1 = 1.0f; + n2 = m.indexOfRefraction; + } + + Fresnel fresnel = calculateFresnel(normal, inRay.direction, n1, n2); + + // Use Russian roulette to determine whether to reflect or refract based on the refractive index + float russianRoulette = xi3; + if (russianRoulette <= fresnel.transmissionCoefficient) + { + // Do refraction + bool internalReflection = false; + outRay.direction = calculateTransmissionDirection(normal, inRay.direction, n1, n2, internalReflection); + color = emittedColor*fresnel.transmissionCoefficient; + if (internalReflection) + return 1; + else + return 2; + } + else + { + // Do reflection + outRay.direction = calculateReflectionDirection(normal, inRay.direction); + color = emittedColor*fresnel.reflectionCoefficient; + return 1; + } + } + else + { + //// Maybe partly specular + //if (m.specularExponent > 0.0f) + //{ + // // Specular reflection + // outRay.direction = calculateReflectionDirection(normal, inRay.direction); + // color = emittedColor; + // return 1; + //} + //// Diffuse surface + + // If surface has reflectivity as well as is diffuse, then run russian rouleete to determine which should be run + float russianRouletteGlossy = xi3; + if (m.reflectivity > 0 && russianRouletteGlossy <= m.reflectivity) + { + // Reflect the ray + outRay.direction = calculateReflectionDirection(normal, inRay.direction); + color = emittedColor; + return 1; + } + + // Run diffuse if we reach here + outRay.direction = calculateRandomDirectionInHemisphere(normal, xi1, xi2, xi3); + color = emittedColor*m.color; + //return 0; + //CookTorrence(inRay, outRay, intersect, normal, emittedColor, color, m, inside, xi1, xi2); + return 0; + } +}; + +#endif \ No newline at end of file diff --git a/src/intersections.h b/src/intersections.h index 714e918..c9b07fb 100755 --- a/src/intersections.h +++ b/src/intersections.h @@ -1,288 +1,674 @@ -// 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 "sceneStructs.h" -#include "cudaMat4.h" -#include "glm/glm.hpp" -#include "utilities.h" -#include - -//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 boxIntersectionTest(glm::vec3 boxMin, glm::vec3 boxMax, staticGeom box, 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)) tymax) || (tymin > tmax) ){ - return -1; - } - if (tymin > tmin){ - tmin = tymin; - } - if (tymax < tmax){ - tmax = tymax; - } - - if((int)rsign.z==0){ - tzmin = (boxMin.z - r.origin.z) * rInverseDirection.z; - tzmax = (boxMax.z - r.origin.z) * rInverseDirection.z; - }else{ - tzmin = (boxMax.z - r.origin.z) * rInverseDirection.z; - tzmax = (boxMin.z - r.origin.z) * rInverseDirection.z; - } - - if ( (tmin > tzmax) || (tzmin > tmax) ){ - return -1; - } - if (tzmin > tmin){ - tmin = tzmin; - } - if (tzmax < tmax){ - tmax = tzmax; - } - if(tmin<0){ - return -1; - } - - glm::vec3 osintersect = r.origin + tmin*r.direction; - - if(abs(osintersect.x-abs(boxMax.x))<.001){ - currentNormal = glm::vec3(1,0,0); - }else if(abs(osintersect.y-abs(boxMax.y))<.001){ - currentNormal = glm::vec3(0,1,0); - }else if(abs(osintersect.z-abs(boxMax.z))<.001){ - currentNormal = glm::vec3(0,0,1); - }else if(abs(osintersect.x+abs(boxMin.x))<.001){ - currentNormal = glm::vec3(-1,0,0); - }else if(abs(osintersect.y+abs(boxMin.y))<.001){ - currentNormal = glm::vec3(0,-1,0); - }else if(abs(osintersect.z+abs(boxMin.z))<.001){ - currentNormal = glm::vec3(0,0,-1); - } - - intersectionPoint = multiplyMV(box.transform, glm::vec4(osintersect, 1.0)); - - - - normal = multiplyMV(box.transform, glm::vec4(currentNormal,0.0)); - return glm::length(intersectionPoint-ro.origin); -} - -//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; - -} - -//Generates a random point on a given sphere -__host__ __device__ glm::vec3 getRandomPointOnSphere(staticGeom sphere, float randomSeed){ - - float radius=.5f; - thrust::default_random_engine rng(hash(randomSeed)); - thrust::uniform_real_distribution u01(0,1); - thrust::uniform_real_distribution u02(-0.5,0.5); - - glm::vec3 point = glm::vec3(0,0,0); - float x=(float)u02(rng); - float y=(float)u02(rng); - float z=0; - float russianRoulette = (float)u01(rng); - if(russianRoulette<0.5){ - z=(float)sqrt(radius*radius-x*x-y*y); - }else - z=-(float)sqrt(radius*radius-x*x-y*y); - - point=glm::vec3(x,y,z); - glm::vec3 randPoint = multiplyMV(sphere.transform, glm::vec4(point,1.0f)); - - return randPoint; -} - +// 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 "sceneStructs.h" +#include "cudaMat4.h" +#include "glm/glm.hpp" +#include "utilities.h" +#include + +//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 boxIntersectionTest(glm::vec3 boxMin, glm::vec3 boxMax, staticGeom box, ray r, glm::vec3& intersectionPoint, glm::vec3& normal); +__host__ __device__ float sphereIntersectionTest(staticGeom sphere, ray r, glm::vec3& intersectionPoint, glm::vec3& normal); +__host__ __device__ double RayPlaneIntersect(glm::vec4 const& P0hom, glm::vec4 const& V0hom, glm::vec3 const& p,glm::vec3 const& normal, glm::mat4 const& T); +__host__ __device__ double RayTriangleIntersect(glm::vec3 const& P0, glm::vec3 const& V0, glm::vec3 const& p1, glm::vec3 const& p2, glm::vec3 const& p3, glm::vec3 normal, glm::mat4 const& T, glm::vec3& surfaceNormal); +__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))::max(); + float tNear = -999999;//std::numeric_limits::min(); + + float t1, t2; + + // For the X planes + if (rt.direction.x == 0) + { + // Ray is || to x-axis + // The light point should be in between the xmin and xmax bounds. Else it doesn't intersect + if (rt.origin.x < xmin || rt.origin.x > xmax) + { + return -1; + } + } + else + { + // T1 = (Xl - Xo) / Xd + t1 = (xmin - rt.origin.x)/rt.direction.x; + + // T2 = (Xh - Xo) / Xd + t2 = (xmax - rt.origin.x)/rt.direction.x; + + // If T1 > T2 swap (T1, T2) /* since T1 intersection with near plane */ + if (t1 > t2) + { + //swap t1 and t2 + double temp = t1; + t1 = t2; + t2 = temp; + } + + // If T1 > Tnear set Tnear =T1 /* want largest Tnear */ + if (t1 > tNear) + tNear = t1; + + // If T2 < Tfar set Tfar="T2" /* want smallest Tfar */ + if (t2 < tFar) + tFar = t2; + + // If Tnear > Tfar box is missed so return false + if (tNear > tFar) + return -1; + + // If Tfar < 0 box is behind ray return false end + if (tFar < 0) + return -1; + } + + // For the Y planes + if (rt.direction.y == 0) + { + // Ray is || to y-axis + // The light point should be in between the ymin and ymax bounds. Else it doesn't intersect + if (rt.origin.y < ymin || rt.origin.y > ymax) + { + return -1; + } + } + else + { + // T1 = (Yl - Yo) / Yd + t1 = (ymin - rt.origin.y)/rt.direction.y; + + // T2 = (Yh - Yo) / Yd + t2 = (ymax - rt.origin.y)/rt.direction.y; + + // If T1 > T2 swap (T1, T2) /* since T1 intersection with near plane */ + if (t1 > t2) + { + //swap t1 and t2 + double temp = t1; + t1 = t2; + t2 = temp; + } + + // If T1 > Tnear set Tnear =T1 /* want largest Tnear */ + if (t1 > tNear) + tNear = t1; + + // If T2 < Tfar set Tfar="T2" /* want smallest Tfar */ + if (t2 < tFar) + tFar = t2; + + // If Tnear > Tfar box is missed so return false + if (tNear > tFar) + return -1; + + // If Tfar < 0 box is behind ray return false end + if (tFar < 0) + return -1; + } + + // For the Z planes + if (rt.direction.z == 0) + { + // Ray is || to z-axis + // The light point should be in between the zmin and zmax bounds. Else it doesn't intersect + if (rt.origin.z < zmin || rt.origin.z > zmax) + { + return -1; + } + } + else + { + // T1 = (Zl - Zo) / Zd + t1 = (zmin - rt.origin.z)/rt.direction.z; + + // T2 = (Zh - Zo) / Zd + t2 = (zmax - rt.origin.z)/rt.direction.z; + + // If T1 > T2 swap (T1, T2) /* since T1 intersection with near plane */ + if (t1 > t2) + { + //swap t1 and t2 + double temp = t1; + t1 = t2; + t2 = temp; + } + + // If T1 > Tnear set Tnear =T1 /* want largest Tnear */ + if (t1 > tNear) + tNear = t1; + + // If T2 < Tfar set Tfar="T2" /* want smallest Tfar */ + if (t2 < tFar) + tFar = t2; + + // If Tnear > Tfar box is missed so return false + if (tNear > tFar) + return -1; + + // If Tfar < 0 box is behind ray return false end + if (tFar < 0) + return -1; + } + + // Box survived all above tests, return with intersection point Tnear and exit point Tfar. + double t; + if (abs(tNear) < 1e-3) + { + if (abs(tFar) < 1e-3) // on the surface + return -1; + t = tFar; + } + else + { + t = tNear; + } + + glm::vec3 p = getPointOnRayUnnormalized(rt, t); + glm::vec4 surNormalTemp = glm::vec4(0.0,0.0,0.0,0.0); + if (p.x <= xmin+(1e-3) && p.x >= xmin-(1e-3)) + surNormalTemp.x = -1; + if (p.y <= ymin+(1e-3) && p.y >= ymin-(1e-3)) + surNormalTemp.y = -1; + if (p.z <= zmin+(1e-3) && p.z >= zmin-(1e-3)) + surNormalTemp.z = -1; + if (p.x <= xmax+(1e-3) && p.x >= xmax-(1e-3)) + surNormalTemp.x = 1; + if (p.y <= ymax+(1e-3) && p.y >= ymax-(1e-3)) + surNormalTemp.y = 1; + if (p.z <= zmax+(1e-3) && p.z >= zmax-(1e-3)) + surNormalTemp.z = 1; + normal = multiplyMV(box.inverseTransposeTransform, surNormalTemp); + normal = glm::normalize(normal); + + intersectionPoint = getPointOnRay(r, t); + //intersectionPoint = multiplyMV(box.transform, glm::vec4(getPointOnRay(rt, t), 1.0)); + return t; + + //glm::vec3 currentNormal = glm::vec3(0,0,0); + + //ray ro = r; + + //glm::vec3 iP0 = multiplyMV(box.inverseTransform,glm::vec4(r.origin, 1.0f)); + //glm::vec3 iP1 = multiplyMV(box.inverseTransform,glm::vec4(r.origin+r.direction, 1.0f)); + //glm::vec3 iV0 = iP1 - iP0; + + //r.origin = iP0; + //r.direction = glm::normalize(iV0); + + //float tmin, tmax, tymin, tymax, tzmin, tzmax; + + //glm::vec3 rsign = getSignOfRay(r); + //glm::vec3 rInverseDirection = getInverseDirectionOfRay(r); + + //if((int)rsign.x==0){ + // tmin = (boxMin.x - r.origin.x) * rInverseDirection.x; + // tmax = (boxMax.x - r.origin.x) * rInverseDirection.x; + //}else{ + // tmin = (boxMax.x - r.origin.x) * rInverseDirection.x; + // tmax = (boxMin.x - r.origin.x) * rInverseDirection.x; + //} + + //if((int)rsign.y==0){ + // tymin = (boxMin.y - r.origin.y) * rInverseDirection.y; + // tymax = (boxMax.y - r.origin.y) * rInverseDirection.y; + //}else{ + // tymin = (boxMax.y - r.origin.y) * rInverseDirection.y; + // tymax = (boxMin.y - r.origin.y) * rInverseDirection.y; + //} + + //if ( (tmin > tymax) || (tymin > tmax) ){ + // return -1; + //} + //if (tymin > tmin){ + // tmin = tymin; + //} + //if (tymax < tmax){ + // tmax = tymax; + //} + + //if((int)rsign.z==0){ + // tzmin = (boxMin.z - r.origin.z) * rInverseDirection.z; + // tzmax = (boxMax.z - r.origin.z) * rInverseDirection.z; + //}else{ + // tzmin = (boxMax.z - r.origin.z) * rInverseDirection.z; + // tzmax = (boxMin.z - r.origin.z) * rInverseDirection.z; + //} + + //if ( (tmin > tzmax) || (tzmin > tmax) ){ + // return -1; + //} + //if (tzmin > tmin){ + // tmin = tzmin; + //} + //if (tzmax < tmax){ + // tmax = tzmax; + //} + //if(tmin<0){ + // return -1; + //} + + //glm::vec3 osintersect = r.origin + tmin*r.direction; + + //if(abs(osintersect.x-abs(boxMax.x))<.001){ + // currentNormal = glm::vec3(1,0,0); + //}else if(abs(osintersect.y-abs(boxMax.y))<.001){ + // currentNormal = glm::vec3(0,1,0); + //}else if(abs(osintersect.z-abs(boxMax.z))<.001){ + // currentNormal = glm::vec3(0,0,1); + //}else if(abs(osintersect.x+abs(boxMin.x))<.001){ + // currentNormal = glm::vec3(-1,0,0); + //}else if(abs(osintersect.y+abs(boxMin.y))<.001){ + // currentNormal = glm::vec3(0,-1,0); + //}else if(abs(osintersect.z+abs(boxMin.z))<.001){ + // currentNormal = glm::vec3(0,0,-1); + //} + + //intersectionPoint = multiplyMV(box.transform, glm::vec4(osintersect, 1.0)); + + + + //normal = multiplyMV(box.transform, glm::vec4(currentNormal,0.0)); + //return glm::length(intersectionPoint-ro.origin); +} + +//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__ double Test_RayPlaneIntersect(glm::vec3 origin, glm::vec3 direction, glm::vec3 const& p, glm::vec3 const& normal) +{ + // If the ray is parallel to the plane and there are either no solutions or an infinite number (line in plane). + // Returning -1 in that case + double dotProd = glm::dot(normal, direction); + if (dotProd == 0.0) + return -1; + + double t = glm::dot(normal, (p-origin)) / dotProd; + + // No intersectionin the viewing direction if t is -ve + if (abs(t) < 1e-3) // on the surface + return 0; + else if (t < 0.0) + return -1; + + return t; +} + +__host__ __device__ double RayTriangleIntersect(staticGeom mesh, ray const& r, glm::vec3 const& p1, glm::vec3 const& p2, glm::vec3 const& p3, + glm::vec3 normal, glm::vec3& intersectionPoint, glm::vec3& surfaceNormal) +{ + //glm::vec4 P0hom = glm::inverse(T) * glm::vec4(P0, 1.0f); + //glm::vec4 V0hom = glm::inverse(T) * glm::vec4(V0, 0.0f); + + glm::vec3 P0hom = multiplyMV(mesh.inverseTransform, glm::vec4(r.origin, 1.0f)); + glm::vec3 V0hom = multiplyMV(mesh.inverseTransform, glm::vec4(glm::normalize(r.direction), 0.0f)); + + //glm::vec4 Rd = glm::vec4(V0hom.x, V0hom.y, V0hom.z, 0.0f); + + if (p1 == p2 || p2 == p3 || p3 == p1) + { + // If any of the points are the same, then return no intersection + return -1; + } + + glm::vec3 s; // s is a point on the plane + if (p1.x == 0.0f && p1.y == 0.0f && p1.z == 0.0f) + s = p2; + else + s = p1; + + //std::vector points; + //points.push_back(p1); + //points.push_back(p2); + //points.push_back(p3); + //glm::vec3 normal = Helper::CalculateNormals(points); + glm::vec4 N = glm::vec4(normal, 0.0f); + surfaceNormal = glm::normalize(multiplyMV(mesh.inverseTransposeTransform, glm::vec4(normal,0.0))); + //surfaceNormal.x = surNormalTemp.x; surfaceNormal.y = surNormalTemp.y; surfaceNormal.z = surNormalTemp.z; + + // If the ray is parallel to the plane and there are either no solutions or an infinite number (line in plane). + // Returning -1 in that case + double dotProd = glm::dot(normal, V0hom); + if (dotProd == 0.0) + return -1; + + double t = Test_RayPlaneIntersect(P0hom, V0hom, s, normal); + + // No intersectionin the viewing direction if t is -ve + if (t <= 0.0) + return -1; + if (abs(t) < 1e-3) // on the surface + return -1; + + glm::vec3 p = P0hom + V0hom*(float)t; + intersectionPoint = r.origin + r.direction*(float)t; + + /*s = area(ΔP1P2P3) + s1 = area(ΔPP2P3) / s + s2 = area(ΔPP3P1) / s + s3 = area(ΔPP1P2) / s*/ + + // Find area of triangle - P1-P2-P3 + glm::mat3 matrix1 = glm::mat3(p1.y, p1.z, 1, + p2.y, p2.z, 1, + p3.y, p3.z, 1); + glm::mat3 matrix2 = glm::mat3(p1.z, p1.x, 1, + p2.z, p2.x, 1, + p3.z, p3.x, 1); + glm::mat3 matrix3 = glm::mat3(p1.x, p1.y, 1, + p2.x, p2.y, 1, + p3.x, p3.y, 1); + double sArea = 0.5*std::sqrt(std::pow(glm::determinant(matrix1),2) + std::pow(glm::determinant(matrix2),2) + std::pow(glm::determinant(matrix3),2)); + + // Find area of triangle - P-P2-P3 + matrix1 = glm::mat3(p.y, p.z, 1, + p2.y, p2.z, 1, + p3.y, p3.z, 1); + matrix2 = glm::mat3(p.z, p.x, 1, + p2.z, p2.x, 1, + p3.z, p3.x, 1); + matrix3 = glm::mat3(p.x, p.y, 1, + p2.x, p2.y, 1, + p3.x, p3.y, 1); + double s1Area = 0.5*std::sqrt(std::pow(glm::determinant(matrix1),2) + std::pow(glm::determinant(matrix2),2) + std::pow(glm::determinant(matrix3),2))/sArea; + + // Find area of triangle - P-P3-P1 + matrix1 = glm::mat3(p.y, p.z, 1, + p3.y, p3.z, 1, + p1.y, p1.z, 1); + matrix2 = glm::mat3(p.z, p.x, 1, + p3.z, p3.x, 1, + p1.z, p1.x, 1); + matrix3 = glm::mat3(p.x, p.y, 1, + p3.x, p3.y, 1, + p1.x, p1.y, 1); + double s2Area = 0.5*std::sqrt(std::pow(glm::determinant(matrix1),2) + std::pow(glm::determinant(matrix2),2) + std::pow(glm::determinant(matrix3),2))/sArea; + + // Find area of triangle - P-P1-P2 + matrix1 = glm::mat3(p.y, p.z, 1, + p1.y, p1.z, 1, + p2.y, p2.z, 1); + matrix2 = glm::mat3(p.z, p.x, 1, + p1.z, p1.x, 1, + p2.z, p2.x, 1); + matrix3 = glm::mat3(p.x, p.y, 1, + p1.x, p1.y, 1, + p2.x, p2.y, 1); + double s3Area = 0.5*std::sqrt(std::pow(glm::determinant(matrix1),2) + std::pow(glm::determinant(matrix2),2) + std::pow(glm::determinant(matrix3),2))/sArea; + + /*P is inside if + 0 ≤ s1 ≤ 1 + 0 ≤ s2 ≤ 1 + 0 ≤ s3 ≤ 1 + s1 + s2 + s3 = 1*/ + if ((s1Area >= 0.0 && s1Area <= 1.0) + && (s2Area >= 0.0 && s2Area <= 1.0) + && (s3Area >= 0.0 && s3Area <= 1.0) + && ((s1Area + s2Area + s3Area) >= 1.0 - 1e-3) + && ((s1Area + s2Area + s3Area) <= 1.0 + 1e-3)) + { + // Point is inside the triangle + return t; + } + + return -1; +} + +// An intersection function calling into each of the individual object intersection functions +__host__ __device__ float findIntersection(staticGeom geom, MeshCuda* mesh, int numberOfMeshes, ray r, glm::vec3& intersectionPoint, glm::vec3& normal){ + switch (geom.type) + { + case GEOMTYPE::SPHERE: + return sphereIntersectionTest(geom, r, intersectionPoint, normal); + case GEOMTYPE::CUBE: + return boxIntersectionTest(geom, r, intersectionPoint, normal); + case GEOMTYPE::MESH: + { + float t = -1; + glm::vec3 temp_intersectionPoint; + glm::vec3 temp_normal; + //printf("NumFaces: %i \n", mesh[geom.meshIndex].numFaces); + for (int currFace = 0; currFace 0.001 && (t < 0 || t_temp < t)) + { + t = t_temp; + intersectionPoint = temp_intersectionPoint; + normal = temp_normal; + } + } + return t; + } + default: + 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; + +} + +//Generates a random point on a given sphere +__host__ __device__ glm::vec3 getRandomPointOnSphere(staticGeom sphere, float randomSeed){ + + float radius=.5f; + thrust::default_random_engine rng(hash(randomSeed)); + thrust::uniform_real_distribution u01(0,1); + thrust::uniform_real_distribution u02(-0.5,0.5); + + glm::vec3 point = glm::vec3(0,0,0); + float x=(float)u02(rng); + float y=(float)u02(rng); + float z=0; + float russianRoulette = (float)u01(rng); + if(russianRoulette<0.5){ + z=(float)sqrt(radius*radius-x*x-y*y); + }else + z=-(float)sqrt(radius*radius-x*x-y*y); + + point=glm::vec3(x,y,z); + glm::vec3 randPoint = multiplyMV(sphere.transform, glm::vec4(point,1.0f)); + + return randPoint; +} + +__host__ __device__ glm::vec3 getRandomPointOnObject(staticGeom geom, float randomSeed) +{ + + switch (geom.type) + { + case GEOMTYPE::SPHERE: + return getRandomPointOnSphere(geom, randomSeed); + case GEOMTYPE::CUBE: + return getRandomPointOnCube(geom, randomSeed); + case GEOMTYPE::MESH: + // TODO - To Be implemented + return glm::vec3(0,0,0); + default: + return glm::vec3(0,0,0); + } +} + #endif \ No newline at end of file diff --git a/src/main.cpp b/src/main.cpp index 4e94892..e74b827 100755 --- a/src/main.cpp +++ b/src/main.cpp @@ -1,398 +1,482 @@ -// 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" - -//------------------------------- -//-------------MAIN-------------- -//------------------------------- - -int main(int argc, char** argv){ - - #ifdef __APPLE__ - // Needed in OSX to force use of OpenGL3.2 - glfwOpenWindowHint(GLFW_OPENGL_VERSION_MAJOR, 3); - glfwOpenWindowHint(GLFW_OPENGL_VERSION_MINOR, 2); - glfwOpenWindowHint(GLFW_OPENGL_FORWARD_COMPAT, GL_TRUE); - glfwOpenWindowHint(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; - } - - // Launch CUDA/GL - - #ifdef __APPLE__ - init(); - #else - init(argc, argv); - #endif - - initCuda(); - - initVAO(); - initTextures(); - - GLuint passthroughProgram; - passthroughProgram = initShader("shaders/passthroughVS.glsl", "shaders/passthroughFS.glsl"); - - glUseProgram(passthroughProgram); - glActiveTexture(GL_TEXTURE0); - - #ifdef __APPLE__ - // send into GLFW main loop - while(1){ - display(); - if (glfwGetKey(GLFW_KEY_ESC) == GLFW_PRESS || !glfwGetWindowParam( GLFW_OPENED )){ - exit(0); - } - } - - glfwTerminate(); - #else - glutDisplayFunc(display); - glutKeyboardFunc(keyboard); - - glutMainLoop(); - #endif - return 0; -} - -//------------------------------- -//---------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(iterationsiterations){ - 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; iobjects.size(); i++){ - geoms[i] = renderScene->objects[i]; - } - for(int i=0; imaterials.size(); i++){ - materials[i] = renderScene->materials[i]; - } - - - // execute the kernel - cudaRaytraceCore(dptr, renderCam, targetFrame, iterations, materials, renderScene->materials.size(), geoms, renderScene->objects.size() ); - - // unmap buffer object - cudaGLUnmapBufferObject(pbo); - }else{ - - if(!finishedRender){ - //output image file - image outputImage(renderCam->resolution.x, renderCam->resolution.y); - - for(int x=0; xresolution.x; x++){ - for(int y=0; yresolution.y; y++){ - int index = x + (y * renderCam->resolution.x); - outputImage.writePixelRGB(x,y,renderCam->image[index]); - } - } - - gammaSettings gamma; - gamma.applyGamma = true; - gamma.gamma = 1.0/2.2; - gamma.divisor = renderCam->iterations; - outputImage.setGammaSettings(gamma); - string filename = renderCam->imageName; - string s; - stringstream out; - out << targetFrame; - s = out.str(); - utilityCore::replaceString(filename, ".bmp", "."+s+".bmp"); - utilityCore::replaceString(filename, ".png", "."+s+".png"); - outputImage.saveImageRGB(filename); - cout << "Saved frame " << s << " to " << filename << endl; - finishedRender = true; - if(singleFrameMode==true){ - cudaDeviceReset(); - exit(0); - } - } - if(targetFrameframes-1){ - - //clear image buffer and move onto next frame - targetFrame++; - iterations = 0; - for(int i=0; iresolution.x*renderCam->resolution.y; i++){ - renderCam->image[i] = glm::vec3(0,0,0); - } - cudaDeviceReset(); - finishedRender = false; - } - } - -} - -#ifdef __APPLE__ - - void display(){ - runCuda(); - - string title = "CIS565 Render | " + utilityCore::convertIntToString(iterations) + " Frames"; - glfwSetWindowTitle(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(); - } - -#else - - void display(){ - runCuda(); - - string title = "565Raytracer | " + utilityCore::convertIntToString(iterations) + " Frames"; - glutSetWindowTitle(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); - - glutPostRedisplay(); - glutSwapBuffers(); - } - - void keyboard(unsigned char key, int x, int y) - { - std::cout << key << std::endl; - switch (key) - { - case(27): - exit(1); - break; - } - } - -#endif - - - - -//------------------------------- -//----------SETUP STUFF---------- -//------------------------------- - -#ifdef __APPLE__ - void init(){ - - if (glfwInit() != GL_TRUE){ - shut_down(1); - } - - // 16 bit color, no depth, alpha or stencil buffers, windowed - if (glfwOpenWindow(width, height, 5, 6, 5, 0, 0, 0, GLFW_WINDOW) != GL_TRUE){ - shut_down(1); - } - - // Set up vertex array object, texture stuff - initVAO(); - initTextures(); - } -#else - void init(int argc, char* argv[]){ - glutInit(&argc, argv); - glutInitDisplayMode(GLUT_DOUBLE | GLUT_RGBA); - glutInitWindowSize(width, height); - glutCreateWindow("565Raytracer"); - - // Init GLEW - glewInit(); - GLenum err = glewInit(); - if (GLEW_OK != err) - { - /* Problem: glewInit failed, something is seriously wrong. */ - std::cout << "glewInit failed, aborting." << std::endl; - exit (1); - } - - initVAO(); - initTextures(); - } -#endif - -void initPBO(GLuint* pbo){ - if (pbo) { - // 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( cutGetMaxGflopsDeviceId() ); - - initPBO(&pbo); - - // Clean up on program exit - atexit(cleanupCuda); - - runCuda(); -} - -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 *vertexShaderPath, const char *fragmentShaderPath){ - GLuint program = glslUtility::createProgram(vertexShaderPath, fragmentShaderPath, attributeLocations, 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; -} - -void shut_down(int return_code){ - #ifdef __APPLE__ - glfwTerminate(); - #endif - exit(return_code); -} +// 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" + +int theButtonState = 0; +int theModifierState = 0; +int lastX = 0, lastY = 0; + +//------------------------------- +//-------------MAIN-------------- +//------------------------------- + +int main(int argc, char** argv){ + + #ifdef __APPLE__ + // Needed in OSX to force use of OpenGL3.2 + glfwOpenWindowHint(GLFW_OPENGL_VERSION_MAJOR, 3); + glfwOpenWindowHint(GLFW_OPENGL_VERSION_MINOR, 2); + glfwOpenWindowHint(GLFW_OPENGL_FORWARD_COMPAT, GL_TRUE); + glfwOpenWindowHint(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; + } + + // Launch CUDA/GL + + #ifdef __APPLE__ + init(); + #else + init(argc, argv); + #endif + + initCuda(); + + initVAO(); + initTextures(); + + GLuint passthroughProgram; + passthroughProgram = initShader("shaders/passthroughVS.glsl", "shaders/passthroughFS.glsl"); + + glUseProgram(passthroughProgram); + glActiveTexture(GL_TEXTURE0); + + #ifdef __APPLE__ + // send into GLFW main loop + while(1){ + display(); + if (glfwGetKey(GLFW_KEY_ESC) == GLFW_PRESS || !glfwGetWindowParam( GLFW_OPENED )){ + exit(0); + } + } + + glfwTerminate(); + #else + glutDisplayFunc(display); + glutKeyboardFunc(keyboard); + glutMouseFunc(onMouseCb); + glutMotionFunc(onMouseMotionCb); + + glutMainLoop(); + #endif + return 0; +} + +//------------------------------- +//---------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(iterationsiterations){ + 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(unsigned int i=0; iobjects.size(); ++i){ + geoms[i] = renderScene->objects[i]; + } + for(unsigned int i=0; imaterials.size(); ++i){ + materials[i] = renderScene->materials[i]; + } + + // execute the kernel + cudaRaytraceCore(dptr, renderCam, targetFrame, iterations, materials, renderScene->materials.size(), + geoms, renderScene->objects.size(), renderScene->meshes); + + // unmap buffer object + cudaGLUnmapBufferObject(pbo); + + // Cleanup + delete[] geoms; + delete[] materials; + }else{ + + if(!finishedRender){ + //output image file + image outputImage(renderCam->resolution.x, renderCam->resolution.y); + + for(int x=0; xresolution.x; x++){ + for(int y=0; yresolution.y; y++){ + int index = x + (y * renderCam->resolution.x); + outputImage.writePixelRGB(renderCam->resolution.x-x,y,renderCam->image[index]); + } + } + + gammaSettings gamma; + gamma.applyGamma = false; + gamma.gamma = 1.0/4.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(targetFrameframes-1){ + + //clear image buffer and move onto next frame + targetFrame++; + iterations = 0; + for(int i=0; iresolution.x*renderCam->resolution.y; i++){ + renderCam->image[i] = glm::vec3(0,0,0); + } + cudaDeviceReset(); + finishedRender = false; + } + } + +} + +#ifdef __APPLE__ + + void display(){ + runCuda(); + + string title = "CIS565 Render | " + utilityCore::convertIntToString(iterations) + " Frames"; + glfwSetWindowTitle(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(); + } + +#else + + void display(){ + runCuda(); + + string title = "565Raytracer | " + utilityCore::convertIntToString(iterations) + " Frames"; + glutSetWindowTitle(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); + + glutPostRedisplay(); + glutSwapBuffers(); + } + + void keyboard(unsigned char key, int x, int y) + { + std::cout << key << std::endl; + switch (key) + { + case(27): + exit(1); + break; + } + } + +#endif + + + + +//------------------------------- +//----------SETUP STUFF---------- +//------------------------------- + +#ifdef __APPLE__ + void init(){ + + if (glfwInit() != GL_TRUE){ + shut_down(1); + } + + // 16 bit color, no depth, alpha or stencil buffers, windowed + if (glfwOpenWindow(width, height, 5, 6, 5, 0, 0, 0, GLFW_WINDOW) != GL_TRUE){ + shut_down(1); + } + + // Set up vertex array object, texture stuff + initVAO(); + initTextures(); + } +#else + void init(int argc, char* argv[]){ + glutInit(&argc, argv); + glutInitDisplayMode(GLUT_DOUBLE | GLUT_RGBA); + glutInitWindowSize(width, height); + glutCreateWindow("565Raytracer"); + + // Init GLEW + glewInit(); + GLenum err = glewInit(); + if (GLEW_OK != err) + { + /* Problem: glewInit failed, something is seriously wrong. */ + std::cout << "glewInit failed, aborting." << std::endl; + exit (1); + } + + initVAO(); + initTextures(); + } +#endif + +void initPBO(GLuint* pbo){ + if (pbo) { + // 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( cutGetMaxGflopsDeviceId() ); + + initPBO(&pbo); + + // Clean up on program exit + atexit(cleanupCuda); + + runCuda(); +} + +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 *vertexShaderPath, const char *fragmentShaderPath){ + GLuint program = glslUtility::createProgram(vertexShaderPath, fragmentShaderPath, attributeLocations, 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; +} + +void shut_down(int return_code){ + #ifdef __APPLE__ + glfwTerminate(); + #endif + exit(return_code); +} + +void onMouseCb(int button, int state, int x, int y) +{ + theButtonState = button; + theModifierState = glutGetModifiers(); + lastX = x; + lastY = y; +} + +void onMouseMotionCb(int x, int y) +{ + int deltaX = lastX - x; + int deltaY = lastY - y; + bool moveLeftRight = abs(deltaX) > abs(deltaY); + bool moveUpDown = !moveLeftRight; + + switch(theButtonState) + { + case GLUT_LEFT_BUTTON: + // Move Camera + if (theModifierState) + { + if (GLUT_ACTIVE_ALT) + { + if (deltaY > 0) + { + renderCam->positions[0].z += 5; + glClear(GL_COLOR_BUFFER_BIT); + iterations = 0; + } + else if (deltaY < 0 && renderCam->positions[0].z > 10) + { + renderCam->positions[0].z -= 5; + glClear(GL_COLOR_BUFFER_BIT); + iterations = 0; + } + } + } + else + { + if (abs(deltaX) > abs(deltaY)) + { + if (deltaX > 0) + { + renderCam->positions[0].x -= 1; + glClear(GL_COLOR_BUFFER_BIT); + iterations = 0; + } + else if (deltaX < 0) + { + renderCam->positions[0].x += 1; + glClear(GL_COLOR_BUFFER_BIT); + iterations = 0; + } + } + else if (deltaY > 0) + { + renderCam->positions[0].y -= 1; + glClear(GL_COLOR_BUFFER_BIT); + iterations = 0; + } + else if (deltaY < 0) + { + renderCam->positions[0].y += 1; + glClear(GL_COLOR_BUFFER_BIT); + iterations = 0; + } + } + break; + } + lastX = x; + lastY = y; + glutPostRedisplay(); +} diff --git a/src/main.h b/src/main.h index 55daf50..9b63adc 100755 --- a/src/main.h +++ b/src/main.h @@ -1,108 +1,110 @@ -// 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 - -#ifndef MAIN_H -#define MAIN_H - -#ifdef __APPLE__ - #include -#else - #include - #include -#endif - -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include "glslUtility.h" -#include "sceneStructs.h" -#include "glm/glm.hpp" -#include "image.h" -#include "raytraceKernel.h" -#include "utilities.h" -#include "scene.h" - -using namespace std; - -//------------------------------- -//----------PATHTRACER----------- -//------------------------------- - -scene* renderScene; -camera* renderCam; -int targetFrame; -int iterations; -bool finishedRender; -bool singleFrameMode; - -//------------------------------- -//------------GL STUFF----------- -//------------------------------- - -GLuint positionLocation = 0; -GLuint texcoordsLocation = 1; -const char *attributeLocations[] = { "Position", "Tex" }; -GLuint pbo = (GLuint)NULL; -GLuint displayImage; - -//------------------------------- -//----------CUDA STUFF----------- -//------------------------------- - -int width=800; int height=800; - -//------------------------------- -//-------------MAIN-------------- -//------------------------------- - -int main(int argc, char** argv); - -//------------------------------- -//---------RUNTIME STUFF--------- -//------------------------------- - -void runCuda(); - -#ifdef __APPLE__ - void display(); -#else - void display(); - void keyboard(unsigned char key, int x, int y); -#endif - -//------------------------------- -//----------SETUP STUFF---------- -//------------------------------- - -#ifdef __APPLE__ - void init(); -#else - void init(int argc, char* argv[]); -#endif - -void initPBO(GLuint* pbo); -void initCuda(); -void initTextures(); -void initVAO(); -GLuint initShader(const char *vertexShaderPath, const char *fragmentShaderPath); - -//------------------------------- -//---------CLEANUP STUFF--------- -//------------------------------- - -void cleanupCuda(); -void deletePBO(GLuint* pbo); -void deleteTexture(GLuint* tex); -void shut_down(int return_code); - +// 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 + +#ifndef MAIN_H +#define MAIN_H + +#ifdef __APPLE__ + #include +#else + #include + #include +#endif + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include "glslUtility.h" +#include "sceneStructs.h" +#include "glm/glm.hpp" +#include "image.h" +#include "raytraceKernel.h" +#include "utilities.h" +#include "scene.h" + +using namespace std; + +//------------------------------- +//----------PATHTRACER----------- +//------------------------------- + +scene* renderScene; +camera* renderCam; +int targetFrame; +unsigned int iterations; +bool finishedRender; +bool singleFrameMode; + +//------------------------------- +//------------GL STUFF----------- +//------------------------------- + +GLuint positionLocation = 0; +GLuint texcoordsLocation = 1; +const char *attributeLocations[] = { "Position", "Tex" }; +GLuint pbo = (GLuint)NULL; +GLuint displayImage; + +//------------------------------- +//----------CUDA STUFF----------- +//------------------------------- + +int width=800; int height=800; + +//------------------------------- +//-------------MAIN-------------- +//------------------------------- + +int main(int argc, char** argv); + +//------------------------------- +//---------RUNTIME STUFF--------- +//------------------------------- + +void runCuda(); + +#ifdef __APPLE__ + void display(); +#else + void display(); + void keyboard(unsigned char key, int x, int y); + void onMouseCb(int button, int state, int x, int y); + void onMouseMotionCb(int x, int y); +#endif + +//------------------------------- +//----------SETUP STUFF---------- +//------------------------------- + +#ifdef __APPLE__ + void init(); +#else + void init(int argc, char* argv[]); +#endif + +void initPBO(GLuint* pbo); +void initCuda(); +void initTextures(); +void initVAO(); +GLuint initShader(const char *vertexShaderPath, const char *fragmentShaderPath); + +//------------------------------- +//---------CLEANUP STUFF--------- +//------------------------------- + +void cleanupCuda(); +void deletePBO(GLuint* pbo); +void deleteTexture(GLuint* tex); +void shut_down(int return_code); + #endif \ No newline at end of file diff --git a/src/raytraceKernel.cu b/src/raytraceKernel.cu index d473c89..e38c842 100755 --- a/src/raytraceKernel.cu +++ b/src/raytraceKernel.cu @@ -1,227 +1,589 @@ -// 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 -#include "utilities.h" -#include "raytraceKernel.h" -#include "intersections.h" -#include "interactions.h" -#include -#include "glm/glm.hpp" - -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)); -} - -//Kernel that does the initial raycast from the camera and caches the result. "First bounce cache, second bounce thrash!" -__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){ - - int index = x + (y * resolution.x); - - thrust::default_random_engine rng(hash(index*time)); - thrust::uniform_real_distribution u01(0,1); - - //standard camera raycast stuff - glm::vec3 E = eye; - glm::vec3 C = view; - glm::vec3 U = up; - float fovx = fov.x; - float fovy = fov.y; - - float CD = glm::length(C); - - glm::vec3 A = glm::cross(C, U); - glm::vec3 B = glm::cross(A, C); - glm::vec3 M = E+C; - glm::vec3 H = (A*float(CD*tan(fovx*(PI/180))))/float(glm::length(A)); - glm::vec3 V = (B*float(CD*tan(-fovy*(PI/180))))/float(glm::length(B)); - - float sx = (x)/(resolution.x-1); - float sy = (y)/(resolution.y-1); - - glm::vec3 P = M + (((2*sx)-1)*H) + (((2*sy)-1)*V); - glm::vec3 PmE = P-E; - glm::vec3 R = E + (float(200)*(PmE))/float(glm::length(PmE)); - - glm::vec3 direction = glm::normalize(R); - //major performance cliff at this point, TODO: find out why! - ray r; - r.origin = eye; - r.direction = direction; - 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, material* materials, int numberOfMaterials){ - - int x = (blockIdx.x * blockDim.x) + threadIdx.x; - int y = (blockIdx.y * blockDim.y) + threadIdx.y; - int index = x + (y * resolution.x); - - ray r = raycastFromCameraKernel(resolution, time, x, y, cam.position, cam.view, cam.up, cam.fov); - - if((x<=resolution.x && y<=resolution.y)){ - - float MAX_DEPTH = 100000000000000000; - float depth = MAX_DEPTH; - - for(int i=0; i-EPSILON){ - MAX_DEPTH = depth; - colors[index] = materials[geoms[i].materialid].color; - } - } - - - - //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, cudamaterials, - numberOfMaterials); - - sendImageToPBO<<>>(PBOpos, renderCam->resolution, cudaimage); - - //retrieve image from GPU - cudaMemcpy( renderCam->image, cudaimage, (int)renderCam->resolution.x*(int)renderCam->resolution.y*sizeof(glm::vec3), cudaMemcpyDeviceToHost); - - //free up stuff, or else we'll leak memory like a madman - cudaFree( cudaimage ); - cudaFree( cudageoms ); - cudaFree( cudamaterials ); - 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 +#include +#include "sceneStructs.h" +#include +#include "utilities.h" +#include "raytraceKernel.h" +#include "intersections.h" +#include "interactions.h" +#include +#include "glm/glm.hpp" + + +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 IsRayAbsent +{ + __host__ __device__ + bool operator()(const RayPackage& rayPackage) + { + return !rayPackage.isPresent; + } +}; + +//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)); +} + +////Kernel that does the initial raycast from the camera and caches the result. "First bounce cache, second bounce thrash!" +//__host__ __device__ ray raycastFromCameraKernel(glm::vec2 resolution, float iterations, float x, float y, glm::vec3 eye, glm::vec3 view, glm::vec3 up, glm::vec2 fov){ +// ray r; +// r.origin = eye; +// +// glm::vec3 A = glm::normalize(glm::cross(view, up)); +// glm::vec3 B = glm::normalize(glm::cross(A, view)); +// +// float tanVert = tan(fov.y*PI/180); +// float tanHor = tan(fov.x*PI/180); +// +// float camDistFromScreen = (float)((resolution.y/2.0)/tanVert); +// glm::vec3 C = view*camDistFromScreen; +// glm::vec3 M = eye + C; +// +// //glm::vec3 H = A * (camDistFromScreen * tanHor); +// //glm::vec3 V = B * (camDistFromScreen * tanVert); +// +// float sx = (x/(float)(resolution.x-1)); +// float sy = (y/(float)(resolution.y-1)); +// glm::vec3 point = M - A*(resolution.x/2.0f)*(2.0f*sx - 1) - B*(resolution.y/2.0f)*(2.0f*sy - 1); +// r.direction = glm::normalize(point - eye); +// +// return r; + + ////standard camera raycast stuff + // glm::vec3 E1 = eye; + // glm::vec3 C = view; + // glm::vec3 U = up; + // float fovx = fov.x; + // float fovy = fov.y; + // + // float CD = glm::length(C); + // + // glm::vec3 A = glm::cross(C, U); + // glm::vec3 B = glm::cross(A, C); + // glm::vec3 M = E1+C; + // glm::vec3 H = (A*float(CD*tan(fovx*(PI/180))))/float(glm::length(A)); + // glm::vec3 V = (B*float(CD*tan(-fovy*(PI/180))))/float(glm::length(B)); + // + // float sx = (x)/(resolution.x-1); + // float sy = (y)/(resolution.y-1); + // + // glm::vec3 P = M + (((2*sx)-1)*H) + (((2*sy)-1)*V); + // glm::vec3 PmE = P-E1; + // glm::vec3 R = E1 + (float(200)*(PmE))/float(glm::length(PmE)); + // + // glm::vec3 direction = glm::normalize(R); + // //major performance cliff at this point, TODO: find out why! + // ray r; + // r.origin = eye; + // r.direction = direction; + // return r; +//} + +__host__ __device__ ray raycastFromCameraKernel(glm::vec2 resolution, float x, float y, glm::vec3 eye, + glm::vec3 M, glm::vec3 A, glm::vec3 B, glm::vec3& point) +{ + ray r; + r.origin = eye; + //glm::vec3 H = A * (camDistFromScreen * tanHor); + //glm::vec3 V = B * (camDistFromScreen * tanVert); + + float sx = (x/(float)(resolution.x-1)); + float sy = (y/(float)(resolution.y-1)); + point = M - A*(resolution.x/2.0f)*(2.0f*sx - 1) - B*(resolution.y/2.0f)*(2.0f*sy - 1); + r.direction = glm::normalize(point - eye); + + return r; +} + +__device__ float GetMinimumIntersection(ray& rayToBeTraced, staticGeom* geoms, unsigned int numberOfGeoms, MeshCuda* meshes, int numberOfMeshes, + unsigned int& minGeomNum, unsigned int& materialId, glm::vec3& minIntersectionPoint, glm::vec3& minNormal) +{ + glm::vec3 intersectionPoint; + glm::vec3 normal; + float t = -1; + + for (unsigned int geomNum = 0; geomNum < numberOfGeoms; ++geomNum) { + float tTemp = findIntersection(geoms[geomNum], meshes, numberOfMeshes, rayToBeTraced, intersectionPoint, normal); + if (tTemp > 0.001f) { + if (t < 0 || tTemp < t) { + t = tTemp; + materialId = geoms[geomNum].materialid; + minIntersectionPoint = intersectionPoint; + minNormal = normal; + minGeomNum = geomNum; + } + } + } + + return t; +} + +//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; + color.y = image[index].y; + color.z = image[index].z; + + float maxColor = color.x; + if (color.y > maxColor) + maxColor = color.y; + if (color.z > maxColor) + maxColor = color.z; + + if (maxColor > 1.0f) + color /= maxColor; + + color *= 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; + } +} + +// When an object intersection has already been found, this function can be called to update RayPackage +// and call the BSDF function +__device__ void FindIntersectionAndBSDF(RayPackage& rayPackage, material mat, glm::vec3 intersectionPoint, glm::vec3 normal, + glm::vec3& color, int randomNumber, int iterations) +{ + // Intersection has occured + + // If the ray has intersected a light source, stop recursion and + // multiply color with light color, else get the color of the material and get the new ray + // Set the color + if (mat.emittance <= 0.001f) + { + // Intersection has occured with non emitting object + thrust::default_random_engine rng(hash(randomNumber)); + thrust::uniform_real_distribution u01(0,1); + + ray temp_r; + temp_r.origin = rayPackage.rayObj.origin; + temp_r.direction = rayPackage.rayObj.direction; + + glm::vec3 currColor = rayPackage.color; + int ret = calculateBSDF(temp_r, rayPackage.rayObj, intersectionPoint, normal, currColor, + rayPackage.color, mat, rayPackage.inside, (float)u01(rng), (float)u01(rng), (float)u01(rng)); + + if (ret == 2) + { + // Transmission. This helps in flipping the refractive indices + rayPackage.inside = !rayPackage.inside; + } + rayPackage.rayObj.origin = intersectionPoint; + rayPackage.isPresent = true; + // The index will remain the same, so no need to change it + } + else + { + color = rayPackage.color * mat.color * mat.emittance; + } +} + +//TODO: IMPLEMENT THIS FUNCTION +//Core raytracer kernel +__global__ void raytraceRay(RayPackage* rayPackage, glm::vec2 resolution, float iterations, int rayDepth, glm::vec3* colors, + staticGeom* geoms, unsigned int numberOfGeoms, material* materials, int numberOfMaterials, MeshCuda* meshes, int numberOfMeshes){ + + int x = (blockIdx.x * blockDim.x) + threadIdx.x; + int y = (blockIdx.y * blockDim.y) + threadIdx.y; + int index = x + (y * resolution.x); + + unsigned int intersectGeomNum; + unsigned int materialId; + glm::vec3 intersectionPoint; + glm::vec3 normal; + + if (rayPackage[index].isPresent) + { + rayPackage[index].isPresent = false; + float t = GetMinimumIntersection(rayPackage[index].rayObj, geoms, numberOfGeoms, meshes, numberOfMeshes, intersectGeomNum, materialId, intersectionPoint, normal); + if (t > 0) + { + FindIntersectionAndBSDF(rayPackage[index], materials[materialId], intersectionPoint, normal, + colors[rayPackage[index].index], rayPackage[index].index*(iterations-1000)*rayDepth, iterations); + } + } +} + +// Kernel call for initial ray cast from the camera +__global__ void InitialRayCast(glm::vec2 resolution, float iterations, cameraData cam, int rayDepth, glm::vec3* colors, + staticGeom* geoms, unsigned int numberOfGeoms, material* materials, unsigned int numberOfMaterials, MeshCuda* meshes, int numberOfMeshes, + RayPackage* rayPackage, glm::vec3 camM, glm::vec3 camA, glm::vec3 camB, float distImagePlaneFromCamera) +{ + int x = (blockIdx.x * blockDim.x) + threadIdx.x; + int y = (blockIdx.y * blockDim.y) + threadIdx.y; + int index = x + (y * resolution.x); + + thrust::default_random_engine rng(hash(index*iterations)); + // Intersection has occured with non emitting object + thrust::uniform_real_distribution u01(-0.5,0.5); + thrust::uniform_real_distribution u02(-0.5,0.5); + float xJittered, yJittered; + xJittered = (float)x + (float)u01(rng); + yJittered = (float)y + (float)u02(rng); + + glm::vec3 screenPoint; + //ray r = raycastFromCameraKernel(resolution, iterations, xJittered, yJittered, cam.position, cam.view, cam.up, cam.fov); + ray r = raycastFromCameraKernel(resolution, xJittered, yJittered, cam.position, camM, camA, camB, screenPoint); + + if (cam.aperture <= 25) + { + // For Depth of field + // Find intersection point of ray with focal plane + // Generate random points within the aperture and shoot it towards the point found above + float focalPlaneDist = cam.focalDist; + glm::vec3 focusPoint; + glm::vec3 rayDirection = r.direction; + + // For intersection with focal plane we can project it in 2D and then find it + // As the focal plane is parallel to the image plane we can do this in 2 projections and find the intersection point + // y coordinate of point + + focusPoint.x = cam.position.x + rayDirection.x * focalPlaneDist; + focusPoint.y = cam.position.y + rayDirection.y * focalPlaneDist; + focusPoint.z = cam.position.z + rayDirection.z * focalPlaneDist; + + //// We know the aperture value (f-Number) from cam.aperture + ////1/f = 1/u + 1/v, u->focalPlane distance from camera, v-> imagePlane distance from camera + //float u = focusPoint.z-cam.position.z; + //float v = cam.position.z - screenPoint.z; + //float focalLength = 1/abs(u) + 1/abs(v); + //focalLength = 1/focalLength; + //if (iterations ==1) + //{ + // printf("Focal Length: %f\n", focalLength); + //} + + //float magnification = focalPlaneDist / distImagePlaneFromCamera; + //float totalDistanceBetweenPlanes = focalPlaneDist + distImagePlaneFromCamera; + + float focalLength = 8.0f; + // f-Number = focalLength/diameter + float lensRadius = (focalLength/cam.aperture)/2.0f; + + // Generate random numbers in the aperture sphere + thrust::uniform_real_distribution u03(-lensRadius, lensRadius); + glm::vec3 rayPointOnCamera = glm::vec3(cam.position.x + (float)u03(rng), cam.position.y + (float)u03(rng), cam.position.z); + r.origin = rayPointOnCamera; + r.direction = glm::normalize(focusPoint - rayPointOnCamera); + + //// Hyperfocal distance H = f*f/(N*c) + //float H = f*f/(cam.aperture*0.1); + //// Depth of focus near point Dnear + //float Dn = (H*focalPlaneDist)/(H+focalPlaneDist); + //float Df = (H*focalPlaneDist)/(H-focalPlaneDist); + //if (iterations == 1 && index <=2) + // std::printf("H: %f Dn: %f Df: %f\n", H, Dn, Df); + + } + + unsigned int intersectGeomNum; + unsigned int materialId; + glm::vec3 intersectionPoint; + glm::vec3 normal; + + // Initialize ray Package + rayPackage[index].rayObj.origin = glm::vec3(0,0,0); + rayPackage[index].rayObj.direction = glm::vec3(0,0,0); + rayPackage[index].inside = false; + rayPackage[index].isPresent = false; + rayPackage[index].color = glm::vec3(1,1,1); + + // Clear pixel value + colors[index] = glm::vec3(0,0,0); + + if((x<=resolution.x && y<=resolution.y)) + { + float t = GetMinimumIntersection(r, geoms, numberOfGeoms, meshes, numberOfMeshes, intersectGeomNum, materialId, intersectionPoint, normal); + if (t > 0) + { + rayPackage[index].rayObj.origin = r.origin; + rayPackage[index].rayObj.direction = r.direction; + + FindIntersectionAndBSDF(rayPackage[index], materials[materialId], intersectionPoint, normal, + colors[index], index*iterations, iterations); + + rayPackage[index].index = index; + } + } +} + +__global__ void CalculateColor(glm::vec2 resolution, glm::vec3* existingColor, glm::vec3* newColor, int 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) + { + newColor[index] = ((existingColor[index] * (float) (iterations-1)) + newColor[index]) / (float)iterations; + } + glm::clamp(newColor[index], 0, 1); +} + + +__global__ void AssignMeshPointers(MeshCuda* mesh, unsigned int meshNumber, glm::vec3* vertices, unsigned int numVertices, + Triangle* faces, unsigned int numFaces, glm::vec3* normals, unsigned int numNormals) +{ + //printf("Old Val: %f \n", vertices[0].x); + (mesh+meshNumber)->vertices = &(vertices[0]); + mesh[meshNumber].faces = faces; + mesh[meshNumber].normals = normals; + mesh[meshNumber].numVertices = numVertices; + mesh[meshNumber].numFaces = numFaces; + mesh[meshNumber].numNormals = numNormals; +} + +__global__ void AssignMeshPointers1(MeshCuda* mesh, int meshNumber, unsigned int p_numVertices, unsigned int p_numNormals, + unsigned int p_numFaces) +{ + mesh[meshNumber].numVertices = p_numVertices; + mesh[meshNumber].numFaces = p_numFaces; + mesh[meshNumber].numNormals = p_numNormals; +} + +__global__ void GetMeshPointers(MeshCuda* mesh, unsigned int meshNumber, glm::vec3** vertices, Triangle** faces, glm::vec3** normals) +{ + vertices = &(mesh[meshNumber].vertices); + faces = &(mesh[meshNumber].faces); + normals = &(mesh[meshNumber].normals); +} + +//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, unsigned int iterations, material* materials, unsigned int numberOfMaterials, + geom* geoms, unsigned int numberOfGeoms, std::vector& meshes){ + + int traceDepth = 1; //determines how many bounces the raytracer traces + int totalRays = (int)renderCam->resolution.x*(int)renderCam->resolution.y; + + // 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)); + + //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; + cam.aperture = renderCam->aperture; + cam.focalDist = renderCam->focalDistance; + + + MeshCuda* cudaMeshPtr = NULL; + int numberOfMeshes = meshes.size(); + MeshCuda* temp_cudaMeshPtr; + if (numberOfMeshes > 0) + { + cudaMalloc((void **)&cudaMeshPtr, numberOfMeshes*sizeof(MeshCuda)); + //temp_cudaMeshPtr = new MeshCuda[numberOfMeshes]; + // Get the meshes into the GPU + for (unsigned int i=0; iresolution, renderCam->positions[frame], renderCam->views[frame], renderCam->ups[frame], + renderCam->fov, M, A, B, distImagePlaneFromCamera); + + //kernel launches + InitialRayCast<<>>(renderCam->resolution, (float)iterations, cam, traceDepth, cudaimage, cudageoms, numberOfGeoms, cudamaterials, + numberOfMaterials, cudaMeshPtr, numberOfMeshes, cudaRayPackage, M, A, B, distImagePlaneFromCamera); + int rayCount = totalRays; + dim3 fullBlocksPerGridNew; + ++traceDepth; + while (traceDepth <= MAX_DEPTH) + { + // wrap raw pointer with a device_ptr for Stream compaction + thrust::device_ptr devRayPackagePtr(cudaRayPackage); + thrust::device_ptr devRayPackageEndPtr = thrust::remove_if(devRayPackagePtr, devRayPackagePtr+rayCount, IsRayAbsent()); + rayCount = devRayPackageEndPtr.get() - devRayPackagePtr.get(); + + if (rayCount <= 0) + break; + + // Create blocks for lesser number of rays found by stream compaction + fullBlocksPerGridNew = dim3((int)ceil(float(renderCam->resolution.x)/float(tileSize)), (int)ceil(float(rayCount)/(float)renderCam->resolution.x/float(tileSize))); + raytraceRay<<>>(cudaRayPackage, renderCam->resolution, (float)iterations, traceDepth, cudaimage, cudageoms, numberOfGeoms, + cudamaterials, numberOfMaterials, cudaMeshPtr, numberOfMeshes); + + ++traceDepth; + } + + glm::vec3* cudaExistingColor = NULL; + cudaMalloc((void**)&cudaExistingColor, totalRays*sizeof(glm::vec3)); + cudaMemcpy( cudaExistingColor, renderCam->image, totalRays*sizeof(glm::vec3), cudaMemcpyHostToDevice); + + CalculateColor<<>>(renderCam->resolution, cudaExistingColor, cudaimage, iterations); + cudaFree( cudaExistingColor ); + + sendImageToPBO<<>>(PBOpos, renderCam->resolution, cudaimage); + + //retrieve image from GPU + cudaMemcpy( renderCam->image, cudaimage, totalRays*sizeof(glm::vec3), cudaMemcpyDeviceToHost); + + //free up stuff, or else we'll leak memory like a madman + cudaFree( cudaimage ); + cudaFree( cudageoms ); + cudaFree( cudamaterials ); + cudaFree( cudaRayPackage ); + delete [] geomList; + + // Clear Mesh Data + for (unsigned int i=0; i>>(cudaMeshPtr, i, &cudaMeshVerticesPtr, &cudaMeshFacesPtr, &cudaMeshNormalsPtr); + + cudaFree(cudaMeshVerticesPtr); + cudaFree(cudaMeshFacesPtr); + cudaFree(cudaMeshNormalsPtr); + } + + if (meshes.size() > 0) + { + cudaFree(cudaMeshPtr); + } + + // make certain the kernel has completed + cudaThreadSynchronize(); + + checkCUDAError("Kernel failed!"); +} diff --git a/src/raytraceKernel.h b/src/raytraceKernel.h index 331e5ce..091121a 100755 --- a/src/raytraceKernel.h +++ b/src/raytraceKernel.h @@ -1,20 +1,21 @@ -// 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 - -#ifndef RAYTRACEKERNEL_H -#define PATHTRACEKERNEL_H - -#include -#include -#include -#include -#include "sceneStructs.h" -#include - -void cudaRaytraceCore(uchar4* pos, camera* renderCam, int frame, int iterations, material* materials, int numberOfMaterials, geom* geoms, int numberOfGeoms); - -#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: +// 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 + +#ifndef RAYTRACEKERNEL_H +#define PATHTRACEKERNEL_H + +#include +#include +#include +#include +#include "sceneStructs.h" +#include + +void cudaRaytraceCore(uchar4* pos, camera* renderCam, int frame, unsigned int iterations, material* materials, unsigned int numberOfMaterials, + geom* geoms, unsigned int numberOfGeoms, std::vector& meshes); + +#endif diff --git a/src/scene.cpp b/src/scene.cpp index f0384b2..13e3bcc 100755 --- a/src/scene.cpp +++ b/src/scene.cpp @@ -1,263 +1,328 @@ -// 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 - -#include -#include "scene.h" - -scene::scene(string filename){ - cout << "Reading scene frome " << filename << "..." << endl; - cout << " " << endl; - char* fname = (char*)filename.c_str(); - fp_in.open(fname); - if(fp_in.is_open()){ - while(fp_in.good()){ - string line; - getline(fp_in, line); - if(!line.empty()){ - vector tokens = utilityCore::tokenizeString(line); - if(strcmp(tokens[0].c_str(), "MATERIAL")==0){ - loadMaterial(tokens[1]); - cout << " " << endl; - }else if(strcmp(tokens[0].c_str(), "OBJECT")==0){ - loadObject(tokens[1]); - cout << " " << endl; - }else if(strcmp(tokens[0].c_str(), "CAMERA")==0){ - loadCamera(); - cout << " " << endl; - } - } - } - } -} - -int scene::loadObject(string objectid){ - int id = atoi(objectid.c_str()); - if(id!=objects.size()){ - cout << "ERROR: OBJECT ID does not match expected number of objects" << endl; - return -1; - }else{ - cout << "Loading Object " << id << "..." << endl; - geom newObject; - string line; - - //load object type - getline(fp_in,line); - if (!line.empty() && fp_in.good()){ - if(strcmp(line.c_str(), "sphere")==0){ - cout << "Creating new sphere..." << endl; - newObject.type = SPHERE; - }else if(strcmp(line.c_str(), "cube")==0){ - 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; - } - } - } - - //link material - getline(fp_in,line); - if(!line.empty() && fp_in.good()){ - vector tokens = utilityCore::tokenizeString(line); - newObject.materialid = atoi(tokens[1].c_str()); - cout << "Connecting Object " << objectid << " to Material " << newObject.materialid << "..." << endl; - } - - //load frames - int frameCount = 0; - getline(fp_in,line); - vector translations; - vector scales; - vector rotations; - while (!line.empty() && fp_in.good()){ - - //check frame number - vector tokens = utilityCore::tokenizeString(line); - if(strcmp(tokens[0].c_str(), "frame")!=0 || atoi(tokens[1].c_str())!=frameCount){ - cout << "ERROR: Incorrect frame count!" << endl; - return -1; - } - - //load tranformations - for(int i=0; i<3; i++){ - glm::vec3 translation; glm::vec3 rotation; glm::vec3 scale; - getline(fp_in,line); - tokens = utilityCore::tokenizeString(line); - if(strcmp(tokens[0].c_str(), "TRANS")==0){ - translations.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(), "ROTAT")==0){ - rotations.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(), "SCALE")==0){ - scales.push_back(glm::vec3(atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str()))); - } - } - - frameCount++; - getline(fp_in,line); - } - - //move frames into CUDA readable arrays - newObject.translations = new glm::vec3[frameCount]; - newObject.rotations = new glm::vec3[frameCount]; - newObject.scales = new glm::vec3[frameCount]; - newObject.transforms = new cudaMat4[frameCount]; - newObject.inverseTransforms = new cudaMat4[frameCount]; - for(int i=0; i tokens = utilityCore::tokenizeString(line); - if(strcmp(tokens[0].c_str(), "RES")==0){ - newCamera.resolution = glm::vec2(atoi(tokens[1].c_str()), atoi(tokens[2].c_str())); - }else if(strcmp(tokens[0].c_str(), "FOVY")==0){ - fovy = atof(tokens[1].c_str()); - }else if(strcmp(tokens[0].c_str(), "ITERATIONS")==0){ - newCamera.iterations = atoi(tokens[1].c_str()); - }else if(strcmp(tokens[0].c_str(), "FILE")==0){ - newCamera.imageName = tokens[1]; - } - } - - //load time variable properties (frames) - int frameCount = 0; - string line; - getline(fp_in,line); - vector positions; - vector views; - vector ups; - while (!line.empty() && fp_in.good()){ - - //check frame number - vector tokens = utilityCore::tokenizeString(line); - if(strcmp(tokens[0].c_str(), "frame")!=0 || atoi(tokens[1].c_str())!=frameCount){ - cout << "ERROR: Incorrect frame count!" << endl; - return -1; - } - - //load camera properties - for(int i=0; i<3; i++){ - //glm::vec3 translation; glm::vec3 rotation; glm::vec3 scale; - getline(fp_in,line); - tokens = utilityCore::tokenizeString(line); - if(strcmp(tokens[0].c_str(), "EYE")==0){ - positions.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(), "VIEW")==0){ - views.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(), "UP")==0){ - ups.push_back(glm::vec3(atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str()))); - } - } - - frameCount++; - getline(fp_in,line); - } - newCamera.frames = frameCount; - - //move frames into CUDA readable arrays - newCamera.positions = new glm::vec3[frameCount]; - newCamera.views = new glm::vec3[frameCount]; - newCamera.ups = new glm::vec3[frameCount]; - for(int i=0; i tokens = utilityCore::tokenizeString(line); - if(strcmp(tokens[0].c_str(), "RGB")==0){ - glm::vec3 color( atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str()) ); - newMaterial.color = color; - }else if(strcmp(tokens[0].c_str(), "SPECEX")==0){ - newMaterial.specularExponent = atof(tokens[1].c_str()); - }else if(strcmp(tokens[0].c_str(), "SPECRGB")==0){ - glm::vec3 specColor( atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str()) ); - newMaterial.specularColor = specColor; - }else if(strcmp(tokens[0].c_str(), "REFL")==0){ - newMaterial.hasReflective = atof(tokens[1].c_str()); - }else if(strcmp(tokens[0].c_str(), "REFR")==0){ - newMaterial.hasRefractive = atof(tokens[1].c_str()); - }else if(strcmp(tokens[0].c_str(), "REFRIOR")==0){ - newMaterial.indexOfRefraction = atof(tokens[1].c_str()); - }else if(strcmp(tokens[0].c_str(), "SCATTER")==0){ - newMaterial.hasScatter = atof(tokens[1].c_str()); - }else if(strcmp(tokens[0].c_str(), "ABSCOEFF")==0){ - glm::vec3 abscoeff( atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str()) ); - newMaterial.absorptionCoefficient = abscoeff; - }else if(strcmp(tokens[0].c_str(), "RSCTCOEFF")==0){ - newMaterial.reducedScatterCoefficient = atof(tokens[1].c_str()); - }else if(strcmp(tokens[0].c_str(), "EMITTANCE")==0){ - newMaterial.emittance = atof(tokens[1].c_str()); - - } - } - materials.push_back(newMaterial); - return 1; - } -} +// 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 + +#include +#include "scene.h" + +scene::scene(string filename){ + cout << "Reading scene frome " << filename << "..." << endl; + cout << " " << endl; + char* fname = (char*)filename.c_str(); + fp_in.open(fname); + if(fp_in.is_open()){ + while(fp_in.good()){ + string line; + getline(fp_in, line); + if(!line.empty()){ + vector tokens = utilityCore::tokenizeString(line); + if(strcmp(tokens[0].c_str(), "MATERIAL")==0){ + loadMaterial(tokens[1]); + cout << " " << endl; + }else if(strcmp(tokens[0].c_str(), "OBJECT")==0){ + loadObject(tokens[1]); + cout << " " << endl; + }else if(strcmp(tokens[0].c_str(), "CAMERA")==0){ + loadCamera(); + cout << " " << endl; + } + } + } + } +} + +int scene::loadObject(string objectid){ + int id = atoi(objectid.c_str()); + if(id!=objects.size()){ + cout << "ERROR: OBJECT ID does not match expected number of objects" << endl; + return -1; + }else{ + cout << "Loading Object " << id << "..." << endl; + geom newObject; + string line; + newObject.meshIndex = -1; + + //load object type + getline(fp_in,line); + if (!line.empty() && fp_in.good()){ + if(strcmp(line.c_str(), "sphere")==0){ + cout << "Creating new sphere..." << endl; + newObject.type = SPHERE; + }else if(strcmp(line.c_str(), "cube")==0){ + 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; + LoadMesh(line); + newObject.meshIndex = meshes.size()-1; + }else{ + cout << "ERROR: " << line << " is not a valid object type!" << endl; + return -1; + } + } + } + + //link material + getline(fp_in,line); + if(!line.empty() && fp_in.good()){ + vector tokens = utilityCore::tokenizeString(line); + newObject.materialid = atoi(tokens[1].c_str()); + cout << "Connecting Object " << objectid << " to Material " << newObject.materialid << "..." << endl; + } + + //load frames + int frameCount = 0; + getline(fp_in,line); + vector translations; + vector scales; + vector rotations; + while (!line.empty() && fp_in.good()){ + + //check frame number + vector tokens = utilityCore::tokenizeString(line); + if(strcmp(tokens[0].c_str(), "frame")!=0 || atoi(tokens[1].c_str())!=frameCount){ + cout << "ERROR: Incorrect frame count!" << endl; + return -1; + } + + //load tranformations + for(int i=0; i<3; i++){ + glm::vec3 translation; glm::vec3 rotation; glm::vec3 scale; + getline(fp_in,line); + tokens = utilityCore::tokenizeString(line); + if(strcmp(tokens[0].c_str(), "TRANS")==0){ + translations.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(), "ROTAT")==0){ + rotations.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(), "SCALE")==0){ + scales.push_back(glm::vec3(atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str()))); + } + } + + frameCount++; + getline(fp_in,line); + } + + //move frames into CUDA readable arrays + newObject.translations = new glm::vec3[frameCount]; + newObject.rotations = new glm::vec3[frameCount]; + newObject.scales = new glm::vec3[frameCount]; + newObject.transforms = new cudaMat4[frameCount]; + newObject.inverseTransforms = new cudaMat4[frameCount]; + newObject.inverseTransposeTransforms = new cudaMat4[frameCount]; + for(int i=0; i tokens = utilityCore::tokenizeString(line); + if(strcmp(tokens[0].c_str(), "RES")==0){ + newCamera.resolution = glm::vec2(atoi(tokens[1].c_str()), atoi(tokens[2].c_str())); + }else if(strcmp(tokens[0].c_str(), "FOVY")==0){ + fovy = atof(tokens[1].c_str()); + }else if(strcmp(tokens[0].c_str(), "ITERATIONS")==0){ + 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(), "APERTURE")==0){ + newCamera.aperture = atof(tokens[1].c_str()); + }else if(strcmp(tokens[0].c_str(), "FOCAL_DISTANCE")==0){ + newCamera.focalDistance = atof(tokens[1].c_str()); + } + } + + //load time variable properties (frames) + int frameCount = 0; + string line; + getline(fp_in,line); + vector positions; + vector views; + vector ups; + while (!line.empty() && fp_in.good()){ + + //check frame number + vector tokens = utilityCore::tokenizeString(line); + if(strcmp(tokens[0].c_str(), "frame")!=0 || atoi(tokens[1].c_str())!=frameCount){ + cout << "ERROR: Incorrect frame count!" << endl; + return -1; + } + + //load camera properties + for(int i=0; i<3; i++){ + //glm::vec3 translation; glm::vec3 rotation; glm::vec3 scale; + getline(fp_in,line); + tokens = utilityCore::tokenizeString(line); + if(strcmp(tokens[0].c_str(), "EYE")==0){ + positions.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(), "VIEW")==0){ + views.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(), "UP")==0){ + ups.push_back(glm::vec3(atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str()))); + } + } + + frameCount++; + getline(fp_in,line); + } + newCamera.frames = frameCount; + + //move frames into CUDA readable arrays + newCamera.positions = new glm::vec3[frameCount]; + newCamera.views = new glm::vec3[frameCount]; + newCamera.ups = new glm::vec3[frameCount]; + for(int i=0; i tokens = utilityCore::tokenizeString(line); + if(strcmp(tokens[0].c_str(), "RGB")==0){ + glm::vec3 color( atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str()) ); + newMaterial.color = color; + }else if(strcmp(tokens[0].c_str(), "SPECEX")==0){ + newMaterial.specularExponent = atof(tokens[1].c_str()); + }else if(strcmp(tokens[0].c_str(), "SPECRGB")==0){ + glm::vec3 specColor( atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str()) ); + newMaterial.specularColor = specColor; + }else if(strcmp(tokens[0].c_str(), "REFL")==0){ + newMaterial.reflectivity = atof(tokens[1].c_str()); + }else if(strcmp(tokens[0].c_str(), "REFR")==0){ + float temp = atof(tokens[1].c_str()); + if (temp > 0.01f) { + newMaterial.hasRefractive = true; + } else { + newMaterial.hasRefractive = false; + } + }else if(strcmp(tokens[0].c_str(), "REFRIOR")==0){ + newMaterial.indexOfRefraction = atof(tokens[1].c_str()); + }else if(strcmp(tokens[0].c_str(), "SCATTER")==0){ + newMaterial.hasScatter = atof(tokens[1].c_str()); + }else if(strcmp(tokens[0].c_str(), "ABSCOEFF")==0){ + glm::vec3 abscoeff( atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str()) ); + newMaterial.absorptionCoefficient = abscoeff; + }else if(strcmp(tokens[0].c_str(), "RSCTCOEFF")==0){ + newMaterial.reducedScatterCoefficient = atof(tokens[1].c_str()); + }else if(strcmp(tokens[0].c_str(), "EMITTANCE")==0){ + newMaterial.emittance = atof(tokens[1].c_str()); + + } + } + materials.push_back(newMaterial); + return 1; + } +} + +int scene::LoadMesh(string filename) +{ + char* fname = (char*)filename.c_str(); + ifstream fp_mesh; + fp_mesh.open(fname); + if(fp_mesh.is_open()) + { + Mesh mesh; + + while(fp_mesh.good()) + { + string line; + getline(fp_mesh, line); + if(!line.empty()) + { + vector tokens = utilityCore::tokenizeString(line); + if(strcmp(tokens[0].c_str(), "v")==0) + { + if (tokens.size() < 4) + { + std::cout << "Error in file data - (vertices)" << std::endl; + return -1; + } + mesh.vertices.push_back(glm::vec3(atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str()))); + } + + if(strcmp(tokens[0].c_str(), "f")==0) + { + if (tokens.size() < 4) + { + std::cout << "Error in file data - (faces)" << std::endl; + return -1; + } + Triangle tri = Triangle(atoi(tokens[1].c_str())-1, atoi(tokens[2].c_str())-1, atoi(tokens[3].c_str())-1); + mesh.faces.push_back(tri); + + // Determine normals + mesh.normals.push_back(mesh.CalculateNormals(mesh.faces.size()-1)); + } + } + } + + if (mesh.vertices.size() > 0) + { + meshes.push_back(mesh); + } + } + + return 0; +} diff --git a/src/scene.h b/src/scene.h index 9bfa71f..7eceb13 100755 --- a/src/scene.h +++ b/src/scene.h @@ -1,34 +1,36 @@ -// 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 SCENE_H -#define SCENE_H - -#include "glm/glm.hpp" -#include "utilities.h" -#include -#include "sceneStructs.h" -#include -#include -#include - -using namespace std; - -class scene{ -private: - ifstream fp_in; - int loadMaterial(string materialid); - int loadObject(string objectid); - int loadCamera(); -public: - scene(string filename); - ~scene(); - - vector objects; - vector materials; - camera renderCam; -}; - -#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 SCENE_H +#define SCENE_H + +#include "glm/glm.hpp" +#include "utilities.h" +#include +#include "sceneStructs.h" +#include +#include +#include + +using namespace std; + +class scene{ +private: + ifstream fp_in; + int loadMaterial(string materialid); + int loadObject(string objectid); + int LoadMesh(string filename); + int loadCamera(); +public: + scene(string filename); + ~scene(); + + vector objects; + vector meshes; + vector materials; + camera renderCam; +}; + +#endif diff --git a/src/sceneStructs.h b/src/sceneStructs.h index b10f1cf..3859dc6 100755 --- a/src/sceneStructs.h +++ b/src/sceneStructs.h @@ -1,76 +1,217 @@ -// 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 CUDASTRUCTS_H -#define CUDASTRUCTS_H - -#include "glm/glm.hpp" -#include "cudaMat4.h" -#include -#include - -enum GEOMTYPE{ SPHERE, CUBE, MESH }; - -struct ray { - glm::vec3 origin; - glm::vec3 direction; -}; - -struct geom { - enum GEOMTYPE type; - int materialid; - int frames; - glm::vec3* translations; - glm::vec3* rotations; - glm::vec3* scales; - cudaMat4* transforms; - cudaMat4* inverseTransforms; -}; - -struct staticGeom { - enum GEOMTYPE type; - int materialid; - glm::vec3 translation; - glm::vec3 rotation; - glm::vec3 scale; - cudaMat4 transform; - cudaMat4 inverseTransform; -}; - -struct cameraData { - glm::vec2 resolution; - glm::vec3 position; - glm::vec3 view; - glm::vec3 up; - glm::vec2 fov; -}; - -struct camera { - glm::vec2 resolution; - glm::vec3* positions; - glm::vec3* views; - glm::vec3* ups; - int frames; - glm::vec2 fov; - unsigned int iterations; - glm::vec3* image; - ray* rayList; - std::string imageName; -}; - -struct material{ - glm::vec3 color; - float specularExponent; - glm::vec3 specularColor; - float hasReflective; - float hasRefractive; - float indexOfRefraction; - float hasScatter; - glm::vec3 absorptionCoefficient; - float reducedScatterCoefficient; - float emittance; -}; - -#endif //CUDASTRUCTS_H +// 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 CUDASTRUCTS_H +#define CUDASTRUCTS_H + +#include "glm/glm.hpp" +#include "cudaMat4.h" +#include +#include +#include + +enum GEOMTYPE{ SPHERE, CUBE, MESH }; + +struct ray { + glm::vec3 origin; + glm::vec3 direction; +}; + +struct Triangle +{ + unsigned int p1; + unsigned int p2; + unsigned int p3; + + Triangle() + { + p1 = 0; + p2 = 0; + p3 = 0; + } + + Triangle(unsigned int m_p1, unsigned int m_p2, unsigned int m_p3) + { + p1 = m_p1; + p2 = m_p2; + p3 = m_p3; + } +}; + +struct Mesh +{ + std::vector vertices; + std::vector faces; + std::vector normals; + + // Calculate the normals for the passed list of vertices of a polygon + // The last vertex will be duplicated as the end vertex, if it is not already present + glm::vec3 CalculateNormals(unsigned int faceIndex) + { + std::vector vertexList; + vertexList.push_back(vertices[faces[faceIndex].p1]); + vertexList.push_back(vertices[faces[faceIndex].p2]); + vertexList.push_back(vertices[faces[faceIndex].p3]); + vertexList.push_back(vertices[faces[faceIndex].p1]); + + float a=0.0f, b=0.0f, c=0.0f; + for(unsigned int i=0; i<(vertexList.size()-1); ++i) + { + a+=((vertexList[i].y - vertexList[i+1].y) * (vertexList[i].z + vertexList[i+1].z)); + b+=((vertexList[i].z - vertexList[i+1].z) * (vertexList[i].x + vertexList[i+1].x)); + c+=((vertexList[i].x - vertexList[i+1].x) * (vertexList[i].y + vertexList[i+1].y)); + } + + return glm::normalize(glm::vec3(a,b,c)); + } +}; + +struct MeshCuda +{ + glm::vec3* vertices; + unsigned int numVertices; + glm::vec3* normals; + unsigned int numNormals; + Triangle* faces; + unsigned int numFaces; + + MeshCuda() + { + vertices = NULL; + faces = NULL; + normals = NULL; + numVertices = 0; + numFaces = 0; + numNormals = 0; + } + + MeshCuda(glm::vec3* p_vertices, unsigned int p_numVertices, glm::vec3* p_normals, unsigned int p_numNormals, + Triangle* p_faces, unsigned int p_numFaces) + { + cudaMalloc((void**)&vertices, p_numVertices*sizeof(glm::vec3)); + cudaMemcpy(vertices, p_vertices, p_numVertices*sizeof(glm::vec3), cudaMemcpyHostToDevice); + + cudaMalloc((void**)&faces, p_numFaces*sizeof(Triangle)); + cudaMemcpy(faces, p_faces, p_numFaces*sizeof(Triangle), cudaMemcpyHostToDevice); + + cudaMalloc((void**)&normals, p_numNormals*sizeof(glm::vec3)); + cudaMemcpy(normals, p_normals, p_numNormals*sizeof(glm::vec3), cudaMemcpyHostToDevice); + + numVertices = p_numVertices; + numFaces = p_numFaces; + numNormals = p_numNormals; + } + + ~MeshCuda() + { + cudaFree(vertices); + cudaFree(faces); + cudaFree(normals); + } +}; + +struct geom { + enum GEOMTYPE type; + int meshIndex; + int materialid; + int frames; + glm::vec3* translations; + glm::vec3* rotations; + glm::vec3* scales; + cudaMat4* transforms; + cudaMat4* inverseTransforms; + cudaMat4* inverseTransposeTransforms; +}; + +struct staticGeom { + enum GEOMTYPE type; + int meshIndex; + int materialid; + glm::vec3 translation; + glm::vec3 rotation; + glm::vec3 scale; + cudaMat4 transform; + cudaMat4 inverseTransform; + cudaMat4 inverseTransposeTransform; +}; + +struct cameraData { + glm::vec2 resolution; + glm::vec3 position; + glm::vec3 view; + glm::vec3 up; + glm::vec2 fov; + float aperture; + float focalDist; +}; + +struct camera { + glm::vec2 resolution; + glm::vec3* positions; + glm::vec3* views; + glm::vec3* ups; + int frames; + glm::vec2 fov; + unsigned int iterations; + glm::vec3* image; + ray* rayList; + std::string imageName; + float aperture; + float focalDistance; +}; + +struct material{ + glm::vec3 color; + float specularExponent; + glm::vec3 specularColor; + float reflectivity; + bool hasRefractive; + float indexOfRefraction; + float hasScatter; + glm::vec3 absorptionCoefficient; + float reducedScatterCoefficient; + float emittance; +}; + +class RayPackage +{ +public: + ray rayObj; + int index; + glm::vec3 color; + bool inside; + bool isPresent; + +public: + RayPackage() + { + rayObj.origin = glm::vec3(0,0,0); + rayObj.direction = glm::vec3(0,0,0); + color = glm::vec3(0,0,0); + inside = false; + isPresent = false; + } + + + RayPackage(ray m_rayObj, glm::vec3 m_color, bool m_inside, bool m_isPresent) + { + rayObj.origin = m_rayObj.origin; + rayObj.direction = m_rayObj.direction; + color = m_color; + inside = m_inside; + isPresent = m_isPresent; + } + + void SetValues(ray m_rayObj, glm::vec3 m_color, bool m_inside, bool m_isPresent) + { + rayObj.origin = m_rayObj.origin; + rayObj.direction = m_rayObj.direction; + color = m_color; + inside = m_inside; + isPresent = m_isPresent; + } +}; + +#endif //CUDASTRUCTS_H diff --git a/src/utilities.h b/src/utilities.h index 5842c33..c6cd151 100755 --- a/src/utilities.h +++ b/src/utilities.h @@ -1,47 +1,48 @@ -// UTILITYCORE- A Utility Library by Yining Karl Li -// This file is part of UTILITYCORE, Coyright (c) 2012 Yining Karl Li -// -// File: utilities.h -// Header for utilities.cpp - -#ifndef Pathtracer_utilities_h -#define Pathtracer_utilities_h - -#include "glm/glm.hpp" -#include -#include -#include -#include -#include -#include -#include -#include "cudaMat4.h" - -const float PI =3.1415926535897932384626422832795028841971; -const float TWO_PI =6.2831853071795864769252867665590057683943; -const float SQRT_OF_ONE_THIRD =0.5773502691896257645091487805019574556476; -const float E =2.7182818284590452353602874713526624977572; -const float EPSILON =.000000001; -const float ZERO_ABSORPTION_EPSILON =0.00001; -const float RAY_BIAS_AMOUNT =0.0002; - -namespace utilityCore { - extern float clamp(float f, float min, float max); - extern bool replaceString(std::string& str, const std::string& from, const std::string& to); - extern glm::vec3 clampRGB(glm::vec3 color); - extern bool epsilonCheck(float a, float b); - extern std::vector tokenizeString(std::string str); - extern cudaMat4 glmMat4ToCudaMat4(glm::mat4 a); - extern glm::mat4 cudaMat4ToGlmMat4(cudaMat4 a); - extern glm::mat4 buildTransformationMatrix(glm::vec3 translation, glm::vec3 rotation, glm::vec3 scale); - extern void printCudaMat4(cudaMat4 m); - extern std::string convertIntToString(int number); - - //----------------------------- - //-------GLM Printers---------- - //----------------------------- - extern void printMat4(glm::mat4); - extern void printVec4(glm::vec4); - extern void printVec3(glm::vec3); -} -#endif +// UTILITYCORE- A Utility Library by Yining Karl Li +// This file is part of UTILITYCORE, Coyright (c) 2012 Yining Karl Li +// +// File: utilities.h +// Header for utilities.cpp + +#ifndef Pathtracer_utilities_h +#define Pathtracer_utilities_h + +#include "glm/glm.hpp" +#include +#include +#include +#include +#include +#include +#include +#include "cudaMat4.h" + +#define PI 3.1415926535897932384626422832795028841971 +#define TWO_PI 6.2831853071795864769252867665590057683943 +#define SQRT_OF_ONE_THIRD 0.5773502691896257645091487805019574556476 +#define E 2.7182818284590452353602874713526624977572 +#define EPSILON .000000001 +#define ZERO_ABSORPTION_EPSILON 0.00001 +#define RAY_BIAS_AMOUNT 0.0002 +#define MAX_DEPTH 10 + +namespace utilityCore { + extern float clamp(float f, float min, float max); + extern bool replaceString(std::string& str, const std::string& from, const std::string& to); + extern glm::vec3 clampRGB(glm::vec3 color); + extern bool epsilonCheck(float a, float b); + extern std::vector tokenizeString(std::string str); + extern cudaMat4 glmMat4ToCudaMat4(glm::mat4 a); + extern glm::mat4 cudaMat4ToGlmMat4(cudaMat4 a); + extern glm::mat4 buildTransformationMatrix(glm::vec3 translation, glm::vec3 rotation, glm::vec3 scale); + extern void printCudaMat4(cudaMat4 m); + extern std::string convertIntToString(int number); + + //----------------------------- + //-------GLM Printers---------- + //----------------------------- + extern void printMat4(glm::mat4); + extern void printVec4(glm::vec4); + extern void printVec3(glm::vec3); +} +#endif