diff --git a/PROJ1_WIN/565Pathtracer.sdf b/PROJ1_WIN/565Pathtracer.sdf new file mode 100644 index 0000000..d4ebc1d Binary files /dev/null and b/PROJ1_WIN/565Pathtracer.sdf differ diff --git a/PROJ1_WIN/565Pathtracer.suo b/PROJ1_WIN/565Pathtracer.suo new file mode 100644 index 0000000..aa05287 Binary files /dev/null and b/PROJ1_WIN/565Pathtracer.suo differ diff --git a/PROJ1_WIN/565Pathtracer/565Pathtracer.vcxproj b/PROJ1_WIN/565Pathtracer/565Pathtracer.vcxproj index 4515c57..e9c8daf 100755 --- a/PROJ1_WIN/565Pathtracer/565Pathtracer.vcxproj +++ b/PROJ1_WIN/565Pathtracer/565Pathtracer.vcxproj @@ -75,7 +75,7 @@ - + @@ -123,6 +123,7 @@ $(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_30,sm_30 @@ -157,7 +158,7 @@ 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) + C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include;C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.2\C\common\inc;../shared/glew/include;../shared/freeglut/include;%(AdditionalIncludeDirectories) Console @@ -170,7 +171,8 @@ $(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 + C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include;C:/ProgramData/NVIDIA Corporation/NVIDIA GPU Computing SDK 4.2/C/common/inc;../shared/glew/includes;../shared/freeglut/includes + compute_30,sm_30 @@ -200,6 +202,6 @@ - + - + \ No newline at end of file diff --git a/PROJ1_WIN/565Pathtracer/565Pathtracer.vcxproj.user b/PROJ1_WIN/565Pathtracer/565Pathtracer.vcxproj.user index dfd9f6c..baf2d1b 100755 --- a/PROJ1_WIN/565Pathtracer/565Pathtracer.vcxproj.user +++ b/PROJ1_WIN/565Pathtracer/565Pathtracer.vcxproj.user @@ -16,4 +16,8 @@ scene=../../scenes/sampleScene.txt WindowsLocalDebugger - + + scene=../../scenes/refractiveGlowingSpheres.txt + WindowsLocalDebugger + + \ No newline at end of file diff --git a/PROJ1_WIN/565Pathtracer/Debug/565Pathtracer.lastbuildstate b/PROJ1_WIN/565Pathtracer/Debug/565Pathtracer.lastbuildstate new file mode 100644 index 0000000..bc1c2bd --- /dev/null +++ b/PROJ1_WIN/565Pathtracer/Debug/565Pathtracer.lastbuildstate @@ -0,0 +1,2 @@ +#v4.0:v100:false +Debug|Win32|C:\Users\ishaan\Documents\penn\565 GPU Progamming\Project2-Pathtracer\PROJ1_WIN\| diff --git a/PROJ1_WIN/565Pathtracer/Debug/565Pathtracer.log b/PROJ1_WIN/565Pathtracer/Debug/565Pathtracer.log new file mode 100644 index 0000000..569d57f --- /dev/null +++ b/PROJ1_WIN/565Pathtracer/Debug/565Pathtracer.log @@ -0,0 +1,41 @@ +Build started 10/14/2013 3:47:59 AM. + 1>Project "C:\Users\ishaan\Documents\penn\565 GPU Progamming\Project2-Pathtracer\PROJ1_WIN\565Pathtracer\565Pathtracer.vcxproj" on node 2 (build target(s)). + 1>PrepareForBuild: + Creating directory "C:\Users\ishaan\Documents\penn\565 GPU Progamming\Project2-Pathtracer\PROJ1_WIN\Debug\". + InitializeBuildStatus: + Creating "Debug\565Pathtracer.unsuccessfulbuild" because "AlwaysCreate" was specified. + AddCudaCompileDeps: + c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\bin\cl.exe /E /nologo /showIncludes /TP /I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.0\include" /I"C:/ProgramData/NVIDIA Corporation/NVIDIA GPU Computing SDK 4.0/C/common/inc" /I../shared/glew/includes /I../shared/freeglut/includes /Ibin /Iinclude /I. /FIcuda_runtime.h /c "C:\Users\ishaan\Documents\penn\565 GPU Progamming\Project2-Pathtracer\src\raytraceKernel.cu" + AddCudaCompilePropsDeps: + Found "CudaToolkitCustomDir". + Found "AdditionalDeps". + Found "CompileOut". + Found "Include". + Found "Keep". + Found "KeepDir". + Found "NvccCompilation". + Found "NvccPath". + Found "RequiredIncludes". + Found "TargetMachinePlatform". + Found "CInterleavedPTX". + Found "CodeGeneration". + Found "GPUDebugInfo". + Found "MaxRegCount". + Found "PtxAsOptionV". + Found "Defines". + Found "Emulation". + Found "HostDebugInfo". + Found "FastMath". + Found "Optimization". + Found "Runtime". + Found "RuntimeChecks". + Found "TypeInfo". + Found "Warning". + Found "AdditionalOptions". + Found "CommandLineTemplate". + 1>C:\Users\ishaan\Documents\penn\565 GPU Progamming\Project2-Pathtracer\PROJ1_WIN\Build\CUDA 4.0.targets(245,9): error : The CUDA Toolkit v4.0 directory '' does not exist. Please verify the CUDA Toolkit is installed properly or define the CudaToolkitDir property to resolve this error. + 1>Done Building Project "C:\Users\ishaan\Documents\penn\565 GPU Progamming\Project2-Pathtracer\PROJ1_WIN\565Pathtracer\565Pathtracer.vcxproj" (build target(s)) -- FAILED. + +Build FAILED. + +Time Elapsed 00:00:00.60 diff --git a/PROJ1_WIN/565Pathtracer/Debug/565Pathtracer.unsuccessfulbuild b/PROJ1_WIN/565Pathtracer/Debug/565Pathtracer.unsuccessfulbuild new file mode 100644 index 0000000..e69de29 diff --git a/PROJ1_WIN/565Pathtracer/Debug/raytraceKernel.cu.cache b/PROJ1_WIN/565Pathtracer/Debug/raytraceKernel.cu.cache new file mode 100644 index 0000000..cdb5d8d --- /dev/null +++ b/PROJ1_WIN/565Pathtracer/Debug/raytraceKernel.cu.cache @@ -0,0 +1,26 @@ +Identity=..\..\src\raytraceKernel.cu +AdditionalDeps= +AdditionalOptions= +CInterleavedPTX=false +CodeGeneration=compute_30,sm_30 +CompileOut=C:\Users\ishaan\Documents\penn\565 GPU Progamming\Project2-Pathtracer\PROJ1_WIN\565Pathtracer\Win32/Debug/raytraceKernel.cu.obj +CudaToolkitCustomDir= +Defines= +Emulation=false +FastMath=false +GPUDebugInfo=true +HostDebugInfo=true +Include=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 +Keep=false +KeepDir=Debug +MaxRegCount=0 +NvccCompilation=compile +NvccPath= +Optimization=Od +PtxAsOptionV=false +RequiredIncludes= +Runtime=MDd +RuntimeChecks=Default +TargetMachinePlatform=32 +TypeInfo=false +Warning=W3 diff --git a/PROJ1_WIN/565Pathtracer/Release/565Pathtracer.Build.CppClean.log b/PROJ1_WIN/565Pathtracer/Release/565Pathtracer.Build.CppClean.log new file mode 100644 index 0000000..31f85eb --- /dev/null +++ b/PROJ1_WIN/565Pathtracer/Release/565Pathtracer.Build.CppClean.log @@ -0,0 +1,21 @@ +C:\USERS\ISHAAN\DOCUMENTS\PENN\565 GPU PROGAMMING\PROJECT2-PATHTRACER\PROJ1_WIN\565PATHTRACER\RELEASE\565PATHTRACER.EXE.INTERMEDIATE.MANIFEST +C:\Users\ishaan\Documents\penn\565 GPU Progamming\Project2-Pathtracer\PROJ1_WIN\565Pathtracer\Release\cl.command.1.tlog +C:\Users\ishaan\Documents\penn\565 GPU Progamming\Project2-Pathtracer\PROJ1_WIN\565Pathtracer\Release\CL.read.1.tlog +C:\Users\ishaan\Documents\penn\565 GPU Progamming\Project2-Pathtracer\PROJ1_WIN\565Pathtracer\Release\CL.write.1.tlog +C:\USERS\ISHAAN\DOCUMENTS\PENN\565 GPU PROGAMMING\PROJECT2-PATHTRACER\PROJ1_WIN\565PATHTRACER\RELEASE\GLSLUTILITY.OBJ +C:\USERS\ISHAAN\DOCUMENTS\PENN\565 GPU PROGAMMING\PROJECT2-PATHTRACER\PROJ1_WIN\565PATHTRACER\RELEASE\IMAGE.OBJ +C:\Users\ishaan\Documents\penn\565 GPU Progamming\Project2-Pathtracer\PROJ1_WIN\565Pathtracer\Release\link.command.1.tlog +C:\Users\ishaan\Documents\penn\565 GPU Progamming\Project2-Pathtracer\PROJ1_WIN\565Pathtracer\Release\link.read.1.tlog +C:\Users\ishaan\Documents\penn\565 GPU Progamming\Project2-Pathtracer\PROJ1_WIN\565Pathtracer\Release\link.write.1.tlog +C:\USERS\ISHAAN\DOCUMENTS\PENN\565 GPU PROGAMMING\PROJECT2-PATHTRACER\PROJ1_WIN\565PATHTRACER\RELEASE\MAIN.OBJ +C:\Users\ishaan\Documents\penn\565 GPU Progamming\Project2-Pathtracer\PROJ1_WIN\565Pathtracer\Release\mt.command.1.tlog +C:\Users\ishaan\Documents\penn\565 GPU Progamming\Project2-Pathtracer\PROJ1_WIN\565Pathtracer\Release\mt.read.1.tlog +C:\Users\ishaan\Documents\penn\565 GPU Progamming\Project2-Pathtracer\PROJ1_WIN\565Pathtracer\Release\mt.write.1.tlog +C:\Users\ishaan\Documents\penn\565 GPU Progamming\Project2-Pathtracer\PROJ1_WIN\565Pathtracer\Release\raytraceKernel.cu.cache +C:\USERS\ISHAAN\DOCUMENTS\PENN\565 GPU PROGAMMING\PROJECT2-PATHTRACER\PROJ1_WIN\565PATHTRACER\RELEASE\SCENE.OBJ +C:\USERS\ISHAAN\DOCUMENTS\PENN\565 GPU PROGAMMING\PROJECT2-PATHTRACER\PROJ1_WIN\565PATHTRACER\RELEASE\STB_IMAGE.OBJ +C:\USERS\ISHAAN\DOCUMENTS\PENN\565 GPU PROGAMMING\PROJECT2-PATHTRACER\PROJ1_WIN\565PATHTRACER\RELEASE\STB_IMAGE_WRITE.OBJ +C:\USERS\ISHAAN\DOCUMENTS\PENN\565 GPU PROGAMMING\PROJECT2-PATHTRACER\PROJ1_WIN\565PATHTRACER\RELEASE\UTILITIES.OBJ +C:\USERS\ISHAAN\DOCUMENTS\PENN\565 GPU PROGAMMING\PROJECT2-PATHTRACER\PROJ1_WIN\565PATHTRACER\RELEASE\VC100.PDB +C:\USERS\ISHAAN\DOCUMENTS\PENN\565 GPU PROGAMMING\PROJECT2-PATHTRACER\PROJ1_WIN\RELEASE\565PATHTRACER.EXE +C:\USERS\ISHAAN\DOCUMENTS\PENN\565 GPU PROGAMMING\PROJECT2-PATHTRACER\PROJ1_WIN\RELEASE\565PATHTRACER.PDB diff --git a/PROJ1_WIN/565Pathtracer/Release/565Pathtracer.log b/PROJ1_WIN/565Pathtracer/Release/565Pathtracer.log new file mode 100644 index 0000000..0e0204f --- /dev/null +++ b/PROJ1_WIN/565Pathtracer/Release/565Pathtracer.log @@ -0,0 +1,15 @@ +Build started 10/14/2013 3:59:17 AM. + 1>Project "C:\Users\ishaan\Documents\penn\565 GPU Progamming\Project2-Pathtracer\PROJ1_WIN\565Pathtracer\565Pathtracer.vcxproj" on node 2 (clean target(s)). + 1>_PrepareForClean: + Deleting file "Release\565Pathtracer.lastbuildstate". + CudaClean: + cmd.exe /C "C:\Users\ishaan\AppData\Local\Temp\5d253cf37e3344ce866e0d5e87c82875.cmd" + "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\bin\nvcc.exe" -ccbin "c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\bin" -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include" -I"C:/ProgramData/NVIDIA Corporation/NVIDIA GPU Computing SDK 4.2/C/common/inc" -I"../shared/glew/includes" -I"../shared/freeglut/includes" -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include" --keep-dir "Release" -maxrregcount=0 --machine 32 --compile -Xcompiler "/EHsc /nologo /Od /Zi /MD " -o "C:\Users\ishaan\Documents\penn\565 GPU Progamming\Project2-Pathtracer\PROJ1_WIN\565Pathtracer\Win32/Release/raytraceKernel.cu.obj" "C:\Users\ishaan\Documents\penn\565 GPU Progamming\Project2-Pathtracer\src\raytraceKernel.cu" -clean + + C:\Users\ishaan\Documents\penn\565 GPU Progamming\Project2-Pathtracer\PROJ1_WIN\565Pathtracer>"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\bin\nvcc.exe" -ccbin "c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\bin" -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include" -I"C:/ProgramData/NVIDIA Corporation/NVIDIA GPU Computing SDK 4.2/C/common/inc" -I"../shared/glew/includes" -I"../shared/freeglut/includes" -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include" --keep-dir "Release" -maxrregcount=0 --machine 32 --compile -Xcompiler "/EHsc /nologo /Od /Zi /MD " -o "C:\Users\ishaan\Documents\penn\565 GPU Progamming\Project2-Pathtracer\PROJ1_WIN\565Pathtracer\Win32/Release/raytraceKernel.cu.obj" "C:\Users\ishaan\Documents\penn\565 GPU Progamming\Project2-Pathtracer\src\raytraceKernel.cu" -clean + Deleting file "Release\..\..\src\raytraceKernel.cu.deps". + 1>Done Building Project "C:\Users\ishaan\Documents\penn\565 GPU Progamming\Project2-Pathtracer\PROJ1_WIN\565Pathtracer\565Pathtracer.vcxproj" (clean target(s)). + +Build succeeded. + +Time Elapsed 00:00:02.87 diff --git a/PROJ1_WIN/565Pathtracer/Release/raytraceKernel.cu.deps b/PROJ1_WIN/565Pathtracer/Release/raytraceKernel.cu.deps new file mode 100644 index 0000000..1d5a64d --- /dev/null +++ b/PROJ1_WIN/565Pathtracer/Release/raytraceKernel.cu.deps @@ -0,0 +1,413 @@ +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\cuda_runtime.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\host_config.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\builtin_types.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\device_types.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\host_defines.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\driver_types.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\host_defines.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\limits.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\crtdefs.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\sal.h +c:\program files (x86)\microsoft visual studio 10.0\vc\include\codeanalysis\sourceannotations.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\vadefs.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\stddef.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\crtdefs.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\surface_types.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\driver_types.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\texture_types.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\driver_types.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\vector_types.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\builtin_types.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\device_types.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\driver_types.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\surface_types.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\texture_types.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\vector_types.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\host_defines.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\channel_descriptor.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\driver_types.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\cuda_runtime_api.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\host_defines.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\builtin_types.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\device_types.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\driver_types.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\surface_types.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\texture_types.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\vector_types.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\host_defines.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\vector_types.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\cuda_runtime_api.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\driver_functions.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\builtin_types.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\device_types.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\driver_types.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\surface_types.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\texture_types.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\vector_types.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\host_defines.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\driver_types.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\host_defines.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\vector_functions.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\builtin_types.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\device_types.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\driver_types.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\surface_types.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\texture_types.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\vector_types.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\host_defines.h +c:\program files\nvidia gpu computing toolkit\cuda\v4.2\include\vector_types.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\stdio.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\crtdefs.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\swprintf.inl +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\cuda.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\stdlib.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\crtdefs.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\cmath +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\yvals.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\crtdefs.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\use_ansi.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\math.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\crtdefs.h +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\sceneStructs.h +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm/glm.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core/_fixes.hpp +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\climits +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\cfloat +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\float.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\crtdefs.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\crtwrn.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\crtdefs.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\limits +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\ymath.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\cwchar +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\wchar.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\crtdefs.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\wtime.inl +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xstddef +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\cstddef +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\cstdio +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core/setup.hpp +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\immintrin.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\wmmintrin.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\nmmintrin.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\smmintrin.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\tmmintrin.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\pmmintrin.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\emmintrin.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xmmintrin.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\mmintrin.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\crtdefs.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\malloc.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\crtdefs.h +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\./core/_detail.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\setup.hpp +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\cassert +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\assert.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\crtdefs.h +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\./core/type.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_half.hpp +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\cstdlib +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_half.inl +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\_detail.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_float.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_half.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\setup.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_int.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\setup.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\_detail.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_gentype.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_size.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_vec1.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_vec.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_gentype.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_float.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_int.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_size.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\_swizzle.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\_swizzle_func.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_vec1.inl +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_vec2.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_vec.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_float.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_int.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_size.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\_swizzle.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_vec2.inl +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_vec3.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_vec.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_float.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_int.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_size.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\_swizzle.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_vec3.inl +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_vec4.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_vec.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_float.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_int.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_size.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\_swizzle.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_vec4.inl +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_mat2x2.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_mat.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_gentype.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_mat2x2.inl +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_mat2x3.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_mat.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_mat2x3.inl +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_mat2x4.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_mat.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_mat2x4.inl +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_mat3x2.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_mat.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_mat3x2.inl +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_mat3x3.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_mat.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_mat3x3.inl +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_mat3x4.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_mat.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_mat3x4.inl +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_mat4x2.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_mat.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_mat4x2.inl +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_mat4x3.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_mat.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_mat4x3.inl +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_mat4x4.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_mat.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\type_mat4x4.inl +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\./core/func_trigonometric.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\func_trigonometric.inl +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\_vectorize.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\./core/func_exponential.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\func_exponential.inl +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\_vectorize.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\./core/func_common.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\_fixes.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\func_common.inl +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\_vectorize.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\./core/func_packing.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\func_packing.inl +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\./core/func_geometric.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\func_geometric.inl +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\_vectorize.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\./core/func_matrix.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\func_matrix.inl +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\_vectorize.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\./core/func_vector_relational.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\_detail.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\func_vector_relational.inl +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\./core/func_integer.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\func_integer.inl +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\_vectorize.hpp +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\intrin.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\crtdefs.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\setjmp.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\crtdefs.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\ammintrin.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\mm3dnow.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\crtdefs.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\mmintrin.h +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\./core/func_noise.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core\func_noise.inl +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\./core/_swizzle.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\cudaMat4.h +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm/glm.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core/_fixes.hpp +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\cuda_runtime.h +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\cuda_runtime.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\string +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\istream +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\ostream +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\ios +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xlocnum +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\streambuf +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xiosbase +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xlocale +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\cstring +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\string.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\crtdefs.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\stdexcept +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\exception +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\eh.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\crtdefs.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xstring +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xmemory +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\new +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xutility +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\utility +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\iosfwd +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\crtdbg.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\crtdefs.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\type_traits +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xtr1common +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxtype_traits +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxtype_traits +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxtype_traits +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxtype_traits +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxtype_traits +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxtype_traits +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxtype_traits +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxtype_traits +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxtype_traits +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxtype_traits +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxtype_traits +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\typeinfo +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xlocinfo +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xlocinfo.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\ctype.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\crtdefs.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\locale.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\crtdefs.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xdebug +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\system_error +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\cerrno +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\errno.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\crtdefs.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\share.h +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm/glm.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core/_fixes.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\utilities.h +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm/glm.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core/_fixes.hpp +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\algorithm +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\memory +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxshared +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxshared +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxshared +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxshared +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxshared +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxshared +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxshared +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxshared +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxshared +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxshared +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxshared +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxshared +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxshared +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxshared +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxshared +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxshared +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxshared +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxshared +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxshared +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxshared +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxshared +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxshared +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxshared +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxshared +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxshared +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxshared +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxshared +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxshared +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxshared +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxshared +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxshared +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxshared +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfwrap1 +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xxshared +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\iterator +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\sstream +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\vector +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\xfunctional +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\cudaMat4.h +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\raytraceKernel.h +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\thrust/random.h +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\thrust/detail/config.h +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\thrust/detail/config/config.h +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\thrust/detail/config/simple_defines.h +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\thrust/detail/config/compiler.h +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\thrust/detail/config/device_backend.h +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\thrust/detail/config/host_device.h +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\host_defines.h +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\thrust/detail/config/debug.h +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\thrust/detail/config/compiler_fence.h +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\thrust/detail/cstdint.h +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\thrust/random/discard_block_engine.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\iostream +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\thrust/random/detail/random_core_access.h +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\thrust/random/detail/discard_block_engine.inl +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\thrust/random/linear_congruential_engine.h +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\thrust/random/detail/linear_congruential_engine_discard.h +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\thrust/random/detail/mod.h +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\thrust/random/detail/linear_congruential_engine.inl +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\thrust/random/linear_feedback_shift_engine.h +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\thrust/random/detail/linear_feedback_shift_engine_wordmask.h +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\thrust/random/detail/linear_feedback_shift_engine.inl +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\thrust/random/subtract_with_carry_engine.h +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\thrust/random/detail/subtract_with_carry_engine.inl +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\thrust/random/xor_combine_engine.h +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\thrust/detail/type_traits.h +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\thrust/detail/type_traits/has_trivial_assign.h +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\thrust/random/detail/xor_combine_engine_max.h +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\thrust/detail/mpl/math.h +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\thrust/random/detail/xor_combine_engine.inl +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\thrust/random/uniform_int_distribution.h +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\thrust/pair.h +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\thrust/detail/pair.inl +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\thrust/detail/integer_traits.h +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\thrust/random/detail/uniform_int_distribution.inl +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\thrust/random/uniform_real_distribution.h +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\thrust/random/detail/uniform_real_distribution.inl +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\thrust/random/normal_distribution.h +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\thrust/random/detail/normal_distribution.inl +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\cuda.h +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\sceneStructs.h +C:/ProgramData/NVIDIA Corporation/NVIDIA GPU Computing SDK 4.2/C/common/inc\cutil_math.h +C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include\cuda_runtime.h +c:\Program Files (x86)\Microsoft Visual Studio 10.0\VC\include\math.h +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\intersections.h +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\sceneStructs.h +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\cudaMat4.h +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm/glm.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\glm\core/_fixes.hpp +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\utilities.h +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\interactions.h +c:\users\ishaan\documents\penn\565 gpu progamming\project2-pathtracer\src\intersections.h +C:/ProgramData/NVIDIA Corporation/NVIDIA GPU Computing SDK 4.2/C/common/inc\cutil_math.h diff --git a/PROJ1_WIN/565Pathtracer/vc100.pdb b/PROJ1_WIN/565Pathtracer/vc100.pdb new file mode 100644 index 0000000..425ddc8 Binary files /dev/null and b/PROJ1_WIN/565Pathtracer/vc100.pdb differ diff --git a/PROJ1_WIN/Release/glew32.dll b/PROJ1_WIN/Release/glew32.dll new file mode 100644 index 0000000..999a59b Binary files /dev/null and b/PROJ1_WIN/Release/glew32.dll differ diff --git a/PROJ1_WIN/ipch/565pathtracer-16d811b2/565pathtracer-9474893a.ipch b/PROJ1_WIN/ipch/565pathtracer-16d811b2/565pathtracer-9474893a.ipch new file mode 100644 index 0000000..7b6ed4d Binary files /dev/null and b/PROJ1_WIN/ipch/565pathtracer-16d811b2/565pathtracer-9474893a.ipch differ diff --git a/PROJ1_WIN/ipch/565pathtracer-16d811b2/565pathtracer-d00238b4.ipch b/PROJ1_WIN/ipch/565pathtracer-16d811b2/565pathtracer-d00238b4.ipch new file mode 100644 index 0000000..00a1282 Binary files /dev/null and b/PROJ1_WIN/ipch/565pathtracer-16d811b2/565pathtracer-d00238b4.ipch differ diff --git a/README.md b/README.md index 324bd3e..37d6829 100755 --- a/README.md +++ b/README.md @@ -3,137 +3,117 @@ CIS565: Project 2: CUDA Pathtracer ------------------------------------------------------------------------------- Fall 2013 ------------------------------------------------------------------------------- -Due Wednesday, 10/02/13 -------------------------------------------------------------------------------- ------------------------------------------------------------------------------- NOTE: ------------------------------------------------------------------------------- -This project requires an NVIDIA graphics card with CUDA capability! Any card after the Geforce 8xxx series will work. If you do not have an NVIDIA graphics card in the machine you are working on, feel free to use any machine in the SIG Lab or in Moore100 labs. All machines in the SIG Lab and Moore100 are equipped with CUDA capable NVIDIA graphics cards. If this too proves to be a problem, please contact Patrick or Liam as soon as possible. - -------------------------------------------------------------------------------- -INTRODUCTION: -------------------------------------------------------------------------------- -In this project, you will extend your raytracer from Project 1 into a full CUDA based global illumination pathtracer. +I'm assuming a Kepler+ generation card with compute capibility 3 upwards. This is for the maximum threads in a block being changed from 512 to 1024 then onwards. This affects my stream compaction. -For this project, you may either choose to continue working off of your codebase from Project 1, or you may choose to use the included basecode in this repository. The basecode for Project 2 is the same as the basecode for Project 1, but with some missing components you will need filled in, such as the intersection testing and camera raycasting methods. - -How you choose to extend your raytracer into a pathtracer is a fairly open-ended problem; the supplied basecode is meant to serve as one possible set of guidelines for doing so, but you may choose any approach you want in your actual implementation, including completely scrapping the provided basecode in favor of your own from-scratch solution. +Also, all test files must be modified to take-in camera aperture and focal distance as parameters. The tags are APERTURE and FOCUSPLANE. Example scenes are there in the scenes folder. ------------------------------------------------------------------------------- -CONTENTS: +INTRODUCTION: ------------------------------------------------------------------------------- -The Project2 root directory contains the following subdirectories: - -* src/ contains the source code for the project. Both the Windows Visual Studio solution and the OSX makefile reference this folder for all source; the base source code compiles on OSX and Windows without modification. -* scenes/ contains an example scene description file. -* renders/ contains two example renders: the raytraced render from Project 1 (GI_no.bmp), and the same scene rendered with global illumination (GI_yes.bmp). -* PROJ1_WIN/ contains a Windows Visual Studio 2010 project and all dependencies needed for building and running on Windows 7. -* PROJ1_OSX/ contains a OSX makefile, run script, and all dependencies needed for building and running on Mac OSX 10.8. -* PROJ1_NIX/ contains a Linux makefile for building and running on Ubuntu - 12.04 LTS. Note that you will need to set the following environment - variables: - - - PATH=$PATH:/usr/local/cuda-5.5/bin - - LD_LIBRARY_PATH=/usr/local/cuda-5.5/lib64:/lib - -The projects build and run exactly the same way as in Project0 and Project1. +I've taken the Ray-Tracer from project 1 and built around to make a path tracer with full global illumination. ------------------------------------------------------------------------------- -REQUIREMENTS: +IMAGES: ------------------------------------------------------------------------------- -In this project, you are given code for: +Lets jump right into the meat here! -* All of the basecode from Project 1, plus: -* Intersection testing code for spheres and cubes -* Code for raycasting from the camera - -You will need to implement the following features. A number of these required features you may have already implemented in Project 1. If you have, you are ahead of the curve and have less work to do! +* Here's some complete images + * Glowing spheres: + ![ScreenShot](renders/glowingSpheres2500_lessMem.png) + * The entire scene is lit wit 6 spheres of low emmission and thus near each sphere there is an area of illumination where it's contribution is maximum. This leads to cool effects while the over all colors are still farily representative of the material colors as the light's have been chosen to add up in a complementary manner to white. Notice the cool fresnel reflections and the caustsics of the lights from the sphere on the top right. + * A variant with the emmisive spheres also being refractive: + ![ScreenShot](renders/refrGlowingSpheres.png) + * Many Spheres: + ![ScreenShot](renders/all2.png) + * This image showcases all the features of the path tracer using a number of spheres. There is a shallow depth of field to focus on the spheres at mid-depth. The multiple caustics arising from not just the light source but the reflections of it as well are a really interesting feature. + -* Full global illumination (including soft shadows, color bleeding, etc.) by pathtracing rays through the scene. -* Properly accumulating emittance and colors to generate a final image -* Supersampled antialiasing -* Parallelization by ray instead of by pixel via stream compaction (you may use Thrust for this). -* Perfect specular reflection +* Diffuse Shading + ![ScreenShot](renders/diffuse.png) -You are also required to implement at least two of the following features. Some of these features you may have already implemented in Project 1. If you have, you may NOT resubmit those features and instead must pick two new ones to implement. +* Specular Shading + ![ScreenShot](renders/specular.0.png) -* From scratch stream compaction (no Thrust). -* Additional BRDF models, such as Cook-Torrance, Ward, etc. Each BRDF model may count as a separate feature. -* Texture mapping -* Bump mapping -* Translational motion blur -* Fresnel-based Refraction, i.e. glass -* OBJ Mesh loading and rendering without KD-Tree -* Interactive camera -* Integrate an existing stackless KD-Tree library, such as CUKD (https://github.com/unvirtual/cukd) -* Depth of field +* Refractive Shading + ![ScreenShot](renders/refr.0.png) -Alternatively, implementing just one of the following features can satisfy the "pick two" feature requirement, since these are correspondingly more difficult problems: +* Fresnel Refraction + ![ScreenShot](renders/fres.0315.png) -* Physically based subsurface scattering and transmission -* Implement and integrate your own stackless KD-Tree from scratch. -* Displacement mapping -* Deformational motion blur +* Fresnel Refraction with Schlick's approximation (Note: There's a small error with the Total internal reflection part of this one which shows up as a black outline on the sphere) + ![ScreenShot](renders/fres_schlick.0350.png) -As yet another alternative, if you have a feature or features you really want to implement that are not on this list, let us know, and we'll probably say yes! +* Painterly Mode + ![ScreenShot](renders/Painted glowing Spheres5000.png) ------------------------------------------------------------------------------- -NOTES ON GLM: +PRE-EXISTING FEATURES: ------------------------------------------------------------------------------- -This project uses GLM, the GL Math library, for linear algebra. You need to know two important points on how GLM is used in this project: - -* In this project, indices in GLM vectors (such as vec3, vec4), are accessed via swizzling. So, instead of v[0], v.x is used, and instead of v[1], v.y is used, and so on and so forth. -* GLM Matrix operations work fine on NVIDIA Fermi cards and later, but pre-Fermi cards do not play nice with GLM matrices. As such, in this project, GLM matrices are replaced with a custom matrix struct, called a cudaMat4, found in cudaMat4.h. A custom function for multiplying glm::vec4s and cudaMat4s is provided as multiplyMV() in intersections.h. +* Bounce Parallel Execution +* Specular Reflection +* Soft shadows and Area Lights +* Depth of Field +* Super Sampled Anti aliasing +* Refraction (albeit not working correctly) +* Interactive camera + * wasdqz for X-Y-Z motion + * [ ] o p for panning camera + * up down arrow keys for changing distance of focus plane + * left right arrow keys for increasing and decreasing the aperture ------------------------------------------------------------------------------- -README +MORE FEATURES: ------------------------------------------------------------------------------- -All students must replace or augment the contents of this Readme.md in a clear -manner with the following: - -* A brief description of the project and the specific features you implemented. -* At least one screenshot of your project running. -* A 30 second or longer video of your project running. To create the video you - can use http://www.microsoft.com/expression/products/Encoder4_Overview.aspx -* A performance evaluation (described in detail below). +The extra features added for this assigment are +* Stream Compaction (My own implementation) +* Fresnel Reflection with Schlick's approximation +* Painterly Mode ------------------------------------------------------------------------------- -PERFORMANCE EVALUATION +OBSERVATIONS: ------------------------------------------------------------------------------- -The performance evaluation is where you will investigate how to make your CUDA -programs more efficient using the skills you've learned in class. You must have -performed at least one experiment on your code to investigate the positive or -negative effects on performance. +Lets take a look at the effects of stream compcation on the visuals: +No Compaction +![ScreenShot](renders/glowingSpheres2500.png) + +With Compaction +![ScreenShot](renders/glowingSpheres2500_lessMem.png) + +The stream compaction strengthens the fresnel reflection visible in the top right sphere. I think this happens because of the "dead" rays not being added at all whatsoever where as in without stream compaction, they would still have had a minute contribution. I'm not sure that's entirely correct though. -One such experiment would be to investigate the performance increase involved -with adding a spatial data-structure to your scene data. +Also, here, I noticed that if I let the code run for 5000 iterations, I started accumulating artifacts. I suspect this is because during the multiplication for the seed of the random number generator, I must have stepped on an upper limit and the bounces for a large iteration all those beyond it would give the same (diffuse) reflected direction. Thus, the artifacts would "add up" into discernable shapes. -Another idea could be looking at the change in timing between various block -sizes. +![ScreenShot](renders/glowingSpheres5000.png) -A good metric to track would be number of rays per second, or frames per -second, or number of objects displayable at 60fps. +An error I encountered early on was for the same seed, I had not been multiplying the same seed with the bounce depth. Thus effectively, all my bounces were in the same direction relative to the local space of the point of intersection. This was leading to light leaks and bright streaks. +![ScreenShot](renders/bloopers/wrongDiffuse.png) +![ScreenShot](renders/bloopers/wrongDiffuse2.png) +![ScreenShot](renders/bloopers/wrongDiffuse3.png) -We encourage you to get creative with your tweaks. Consider places in your code -that could be considered bottlenecks and try to improve them. +Thus, I tried playing around with the seed and realized that upon giving it a constant value, it appeared that the light and all reflections of it seemed to be projected on to all the surfaces. +![ScreenShot](renders/bloopers/non-randomness.png) -Each student should provide no more than a one page summary of their -optimizations along with tables and or graphs to visually explain any -performance differences. +This when let run for a large number of iterations adds up to a very interesting effect where it looks like the scene has been painted with brush strokes yet has correct global illumination properties (soft shadows, color bleeds and caustics). This looked really cool to me so I decided to support it as a "feature" in the code and can be switched on and off! +![ScreenShot](renders/bloopers/paint.png) +![ScreenShot](renders/bloopers/Paint2.png) ------------------------------------------------------------------------------- -THIRD PARTY CODE POLICY +PERFORMANCE ANALYSIS: ------------------------------------------------------------------------------- -* Use of any third-party code must be approved by asking on the Google group. If it is approved, all students are welcome to use it. Generally, we approve use of third-party code that is not a core part of the project. For example, for the ray tracer, we would approve using a third-party library for loading models, but would not approve copying and pasting a CUDA function for doing refraction. -* Third-party code must be credited in README.md. -* Using third-party code without its approval, including using another student's code, is an academic integrity violation, and will result in you receiving an F for the semester. +For performance, I tried two things: +* Stream Compaction: I wanted to see the effects of stream compaction and so the scene has one of the 6 faces of the room empty (behind the camera). I tested stream compaction on two scenes: The glowing spheres scene and the scene with all my features working on multiple spheres. The graphs below plot the average render time per frame vs the number of bounces. It is easy to see that for a small bounce depth the over head of stream compaction is not worth it. Nonetheless, that overhead becomes extremely small compared to the wasted rays being bounced around for a relatively small bounce depth itself! And on the higher side of things, you can see an almost 20x speed up for 10,000 bounce depths. + * ![ScreenShot](renders/glowingSpheresTime.PNG) + * ![ScreenShot](renders/manySpheresTime.PNG) +* Global Memory usage: I had written a lot of my code without storing data from global into local registers and so each time this data was read, it would invoke a global memory call. I removed as many of these as I could and I expected to see a dramatic drop in the run time but there wasn't any noticeable difference. I speculate the NVidia CUDA compiler is smart enough to do these replacements at compilation time itself thus making this exercise futile for the speed up I was hoping to gain but helps decrypt the CUDA compiler a little bit. ------------------------------------------------------------------------------- -SELF-GRADING +THIRD PARTY CODE POLICY ------------------------------------------------------------------------------- -* On the submission date, email your grade, on a scale of 0 to 100, to Liam, liamboone+cis565@gmail.com, with a one paragraph explanation. Be concise and realistic. Recall that we reserve 30 points as a sanity check to adjust your grade. Your actual grade will be (0.7 * your grade) + (0.3 * our grade). We hope to only use this in extreme cases when your grade does not realistically reflect your work - it is either too high or too low. In most cases, we plan to give you the exact grade you suggest. -* Projects are not weighted evenly, e.g., Project 0 doesn't count as much as the path tracer. We will determine the weighting at the end of the semester based on the size of each project. +There exists a shared memory scan function taken from Ch 39 of GPU Gems but it has not been connected into the code. I wanted to use it as a base for understanding the shared memory stream compaction but ended up using my own function based on a naive shared scan instead. ------------------------------------------------------------------------------- SUBMISSION diff --git a/renders/Fresnel.png b/renders/Fresnel.png new file mode 100644 index 0000000..e5b0a56 Binary files /dev/null and b/renders/Fresnel.png differ diff --git a/renders/Painted glowing Spheres5000.png b/renders/Painted glowing Spheres5000.png new file mode 100644 index 0000000..d8d5789 Binary files /dev/null and b/renders/Painted glowing Spheres5000.png differ diff --git a/renders/SphereGlow.0.png b/renders/SphereGlow.0.png new file mode 100644 index 0000000..4223c41 Binary files /dev/null and b/renders/SphereGlow.0.png differ diff --git a/renders/all.png b/renders/all.png new file mode 100644 index 0000000..9c6f841 Binary files /dev/null and b/renders/all.png differ diff --git a/renders/all2.png b/renders/all2.png new file mode 100644 index 0000000..4eb6c57 Binary files /dev/null and b/renders/all2.png differ diff --git a/renders/bloopers/Paint2.png b/renders/bloopers/Paint2.png new file mode 100644 index 0000000..15c715e Binary files /dev/null and b/renders/bloopers/Paint2.png differ diff --git a/renders/bloopers/non-randomness.png b/renders/bloopers/non-randomness.png new file mode 100644 index 0000000..bbcc5f3 Binary files /dev/null and b/renders/bloopers/non-randomness.png differ diff --git a/renders/bloopers/paint.png b/renders/bloopers/paint.png new file mode 100644 index 0000000..d32bc29 Binary files /dev/null and b/renders/bloopers/paint.png differ diff --git a/renders/bloopers/wrongDiffuse.png b/renders/bloopers/wrongDiffuse.png new file mode 100644 index 0000000..ce82134 Binary files /dev/null and b/renders/bloopers/wrongDiffuse.png differ diff --git a/renders/bloopers/wrongDiffuse2.png b/renders/bloopers/wrongDiffuse2.png new file mode 100644 index 0000000..5878c9b Binary files /dev/null and b/renders/bloopers/wrongDiffuse2.png differ diff --git a/renders/bloopers/wrongDiffuse3.png b/renders/bloopers/wrongDiffuse3.png new file mode 100644 index 0000000..35b753e Binary files /dev/null and b/renders/bloopers/wrongDiffuse3.png differ diff --git a/renders/diffuse.png b/renders/diffuse.png new file mode 100644 index 0000000..59ec7b2 Binary files /dev/null and b/renders/diffuse.png differ diff --git a/renders/fres.0315.png b/renders/fres.0315.png new file mode 100644 index 0000000..50a3341 Binary files /dev/null and b/renders/fres.0315.png differ diff --git a/renders/fres_schlick.0350.png b/renders/fres_schlick.0350.png new file mode 100644 index 0000000..9a9272d Binary files /dev/null and b/renders/fres_schlick.0350.png differ diff --git a/renders/glowingSpheres2500.png b/renders/glowingSpheres2500.png new file mode 100644 index 0000000..24be53c Binary files /dev/null and b/renders/glowingSpheres2500.png differ diff --git a/renders/glowingSpheres2500_lessMem.png b/renders/glowingSpheres2500_lessMem.png new file mode 100644 index 0000000..da78ebe Binary files /dev/null and b/renders/glowingSpheres2500_lessMem.png differ diff --git a/renders/glowingSpheres5000.png b/renders/glowingSpheres5000.png new file mode 100644 index 0000000..0ae2587 Binary files /dev/null and b/renders/glowingSpheres5000.png differ diff --git a/renders/glowingSpheresTime.PNG b/renders/glowingSpheresTime.PNG new file mode 100644 index 0000000..4a82dab Binary files /dev/null and b/renders/glowingSpheresTime.PNG differ diff --git a/renders/manySpheresTime.PNG b/renders/manySpheresTime.PNG new file mode 100644 index 0000000..a387f6c Binary files /dev/null and b/renders/manySpheresTime.PNG differ diff --git a/renders/refr.0.png b/renders/refr.0.png new file mode 100644 index 0000000..85fad53 Binary files /dev/null and b/renders/refr.0.png differ diff --git a/renders/refrGlowingSpheres.png b/renders/refrGlowingSpheres.png new file mode 100644 index 0000000..546f42b Binary files /dev/null and b/renders/refrGlowingSpheres.png differ diff --git a/renders/specular.0.png b/renders/specular.0.png new file mode 100644 index 0000000..f4f9a86 Binary files /dev/null and b/renders/specular.0.png differ diff --git a/scenes/diffuse.txt b/scenes/diffuse.txt new file mode 100644 index 0000000..706048d --- /dev/null +++ b/scenes/diffuse.txt @@ -0,0 +1,184 @@ +MATERIAL 0 //white diffuse +RGB 1 1 1 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 0 +REFRIOR 0 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 0 + +MATERIAL 1 //red diffuse +RGB .63 .06 .04 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 0 +REFRIOR 0 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 0 + +MATERIAL 2 //green diffuse +RGB .15 .48 .09 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 0 +REFRIOR 0 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 0 + +MATERIAL 3 //red glossy +RGB .63 .06 .04 +SPECEX 0 +SPECRGB 1 1 1 +REFL 1 +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 1 +REFR 0 +REFRIOR 2 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 0 + +MATERIAL 5 //glass +RGB 1 1 1 +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 .15 .48 .09 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 0 +REFRIOR 2.6 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 0 + +MATERIAL 7 //light +RGB 0.25 0.65 0.85 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 10 + +MATERIAL 8 //light +RGB 1 1 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 15 + +CAMERA +RES 800 800 +FOVY 35 +ITERATIONS 1500 +FILE specular.png +frame 0 +EYE 0 5 12 +VIEW 0 0 -1 +UP 0 1 0 +APERTURE 0 +FOCUSPLANE 12.5 + +OBJECT 0 +cube +material 0 +frame 0 +TRANS 0 0 0 +ROTAT 0 0 90 +SCALE .01 10 10 + +OBJECT 1 +cube +material 0 +frame 0 +TRANS 0 5 -5 +ROTAT 0 90 0 +SCALE .01 10 10 + +OBJECT 2 +cube +material 0 +frame 0 +TRANS 0 10 0 +ROTAT 0 0 90 +SCALE .01 10 10 + +OBJECT 3 +cube +material 1 +frame 0 +TRANS -5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +OBJECT 4 +cube +material 2 +frame 0 +TRANS 5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +OBJECT 5 +sphere +material 5 +frame 0 +TRANS 2 4 0 +ROTAT 35 45 0 +SCALE 3 3 3 + +OBJECT 6 +cube +material 8 +frame 0 +TRANS 0 10 0 +ROTAT 0 0 90 +SCALE .3 3 3 + +OBJECT 7 +sphere +material 7 +frame 0 +TRANS 5 8 0 +ROTAT 0 0 0 +SCALE 1 1 1 + diff --git a/scenes/refractiveGlowingSpheres.txt b/scenes/refractiveGlowingSpheres.txt new file mode 100644 index 0000000..e2c4fa0 --- /dev/null +++ b/scenes/refractiveGlowingSpheres.txt @@ -0,0 +1,303 @@ +MATERIAL 0 //white diffuse +RGB 1 1 1 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 0 +REFRIOR 0 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 0 + +MATERIAL 1 //red diffuse +RGB .63 .06 .04 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 0 +REFRIOR 0 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 0 + +MATERIAL 2 //green diffuse +RGB .15 .48 .09 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 0 +REFRIOR 0 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 0 + +MATERIAL 3 //red glossy +RGB .63 .06 .04 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 0 +REFRIOR 2 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 0 + +MATERIAL 4 //white glossy +RGB 1 1 1 +SPECEX 0 +SPECRGB 1 1 1 +REFL 1 +REFR 0 +REFRIOR 2 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 0 + +MATERIAL 5 //glass +RGB 1 1 1 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 1 +REFRIOR 1.3 +SCATTER 0 +ABSCOEFF .02 5.1 5.7 +RSCTCOEFF 13 +EMITTANCE 0 + +MATERIAL 6 //green glossy +RGB .15 .48 .09 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 0 +REFRIOR 2.6 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 0 + +MATERIAL 7 //light +RGB 1 1 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 1 + +MATERIAL 8 //light +RGB 1 1 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 15 + +MATERIAL 9 // magentalight +RGB 1 0 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 1.55 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 3 + +MATERIAL 10 // bluelight +RGB 0 0 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 1.55 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 3 + +MATERIAL 11 // redlight +RGB 1 0 0 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 1.55 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 3 + +MATERIAL 12 // yellowlight +RGB 1 1 0 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 1.55 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 3 + +MATERIAL 13 // greenlight +RGB 0 1 0 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 1.55 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 3 + +MATERIAL 14 //cyanlight +RGB 0 1 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 1 +REFRIOR 1.55 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 3 + +CAMERA +RES 800 800 +FOVY 35 +ITERATIONS 250 +FILE SphereGlow.png +frame 0 +EYE 0 5 12 +VIEW 0 0 -1 +UP 0 1 0 +APERTURE 0 +FOCUSPLANE 10 + +OBJECT 0 +cube +material 0 +frame 0 +TRANS 0 0 0 +ROTAT 0 0 90 +SCALE .1 10 10 + +OBJECT 1 +cube +material 0 +frame 0 +TRANS 0 5 -5 +ROTAT 0 90 0 +SCALE .1 10 10 + +OBJECT 2 +cube +material 0 +frame 0 +TRANS 0 10 0 +ROTAT 0 0 90 +SCALE .1 10 10 + +OBJECT 3 +cube +material 1 +frame 0 +TRANS -5 5 0 +ROTAT 0 0 0 +SCALE .1 10 10 + +OBJECT 4 +cube +material 2 +frame 0 +TRANS 5 5 0 +ROTAT 0 0 0 +SCALE .1 10 10 + +OBJECT 5 +sphere +material 4 +frame 0 +TRANS 2 8 0 +ROTAT 45 45 0 +SCALE 2 2 2 + +OBJECT 6 +sphere +material 13 +frame 0 +TRANS 0 2 4 +ROTAT 0 0 0 +SCALE 1.5 1.5 1.5 + +OBJECT 7 +sphere +material 12 +frame 0 +TRANS 3 3 1 +ROTAT 0 0 0 +SCALE 1.5 1.5 1.5 + +OBJECT 8 +sphere +material 14 +frame 0 +TRANS -3 3 1 +ROTAT 0 0 0 +SCALE 1.5 1.5 1.5 + +OBJECT 9 +sphere +material 11 +frame 0 +TRANS 3 6 -1 +ROTAT 0 0 0 +SCALE 1.5 1.5 1.5 + +OBJECT 10 +sphere +material 10 +frame 0 +TRANS -3 6 -1 +ROTAT 0 0 0 +SCALE 1.5 1.5 1.5 + +OBJECT 11 +sphere +material 9 +frame 0 +TRANS 0 8 -4 +ROTAT 0 0 0 +SCALE 1.5 1.5 1.5 + +OBJECT 12 +cube +material 0 +frame 0 +TRANS 0 5 0 +ROTAT 45 45 0 +SCALE 2 2 2 + +OBJECT 13 +sphere +material 5 +frame 0 +TRANS -2 8 0 +ROTAT 45 45 0 +SCALE 2 2 2 \ No newline at end of file diff --git a/scenes/sampleScene.txt b/scenes/sampleScene.txt index 52d079e..691f01b 100755 --- a/scenes/sampleScene.txt +++ b/scenes/sampleScene.txt @@ -1,5 +1,5 @@ MATERIAL 0 //white diffuse -RGB 0.9 0.9 0.9 +RGB 1 1 1 SPECEX 0 SPECRGB 1 1 1 REFL 0 @@ -35,10 +35,10 @@ RSCTCOEFF 0 EMITTANCE 0 MATERIAL 3 //red glossy -RGB .63 .26 .24 +RGB .63 .06 .04 SPECEX 0 SPECRGB 1 1 1 -REFL 0 +REFL 1 REFR 0 REFRIOR 2 SCATTER 0 @@ -50,7 +50,7 @@ MATERIAL 4 //white glossy RGB 1 1 1 SPECEX 0 SPECRGB 1 1 1 -REFL 0 +REFL 1 REFR 0 REFRIOR 2 SCATTER 0 @@ -59,19 +59,19 @@ RSCTCOEFF 0 EMITTANCE 0 MATERIAL 5 //glass -RGB 0 0 0 +RGB 1 1 1 SPECEX 0 SPECRGB 1 1 1 REFL 0 REFR 1 -REFRIOR 2.2 +REFRIOR 1.5 SCATTER 0 ABSCOEFF .02 5.1 5.7 RSCTCOEFF 13 EMITTANCE 0 MATERIAL 6 //green glossy -RGB .35 .48 .29 +RGB .15 .48 .09 SPECEX 0 SPECRGB 1 1 1 REFL 0 @@ -108,17 +108,15 @@ EMITTANCE 15 CAMERA RES 800 800 -FOVY 25 -ITERATIONS 5000 -FILE test.bmp +FOVY 35 +ITERATIONS 2500 +FILE test.png frame 0 EYE 0 4.5 12 VIEW 0 0 -1 UP 0 1 0 -frame 1 -EYE 0 4.5 12 -VIEW 0 0 -1 -UP 0 1 0 +APERTURE 1.75 +FOCUSPLANE 12.5 OBJECT 0 cube @@ -127,10 +125,6 @@ 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 @@ -139,10 +133,6 @@ 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 @@ -151,10 +141,6 @@ 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 @@ -163,10 +149,6 @@ 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 @@ -175,10 +157,6 @@ 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 @@ -187,10 +165,6 @@ frame 0 TRANS 0 2 0 ROTAT 0 180 0 SCALE 3 3 3 -frame 1 -TRANS 0 2 0 -ROTAT 0 180 0 -SCALE 3 3 3 OBJECT 6 sphere @@ -199,10 +173,6 @@ frame 0 TRANS 2 5 2 ROTAT 0 180 0 SCALE 2.5 2.5 2.5 -frame 1 -TRANS 2 5 2 -ROTAT 0 180 0 -SCALE 2.5 2.5 2.5 OBJECT 7 sphere @@ -211,19 +181,29 @@ frame 0 TRANS -2 5 -2 ROTAT 0 180 0 SCALE 3 3 3 -frame 1 -TRANS -2 5 -2 -ROTAT 0 180 0 -SCALE 3 3 3 + OBJECT 8 cube -material 8 +material 8 frame 0 TRANS 0 10 0 ROTAT 0 0 90 SCALE .3 3 3 -frame 1 -TRANS 0 10 0 -ROTAT 0 0 90 -SCALE .3 3 3 + +OBJECT 9 +sphere +material 4 +frame 0 +TRANS 3 8 0 +ROTAT 0 0 0 +SCALE 2 2 2 + + +OBJECT 10 +sphere +material 5 +frame 0 +TRANS -3 3 0 +ROTAT 0 45 0 +SCALE 3 3 3 diff --git a/scenes/sampleScene_refr.txt b/scenes/sampleScene_refr.txt new file mode 100644 index 0000000..863f8d8 --- /dev/null +++ b/scenes/sampleScene_refr.txt @@ -0,0 +1,208 @@ +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.3 +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.8 +REFR 0 +REFRIOR 0 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 0 + +MATERIAL 3 //red glossy +RGB .63 .06 .04 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0.2 +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.2 +REFR 0 +REFRIOR 2 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 0 + +MATERIAL 5 //glass +RGB 1 1 1 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 1 +REFRIOR 2.2 +SCATTER 0 +ABSCOEFF .02 5.1 5.7 +RSCTCOEFF 13 +EMITTANCE 0 + +MATERIAL 6 //green glossy +RGB .15 .48 .09 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0.1 +REFR 0 +REFRIOR 2.6 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 0 + +MATERIAL 7 //light +RGB 1 1 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 1 + +MATERIAL 8 //light +RGB 1 1 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 15 + +CAMERA +RES 800 800 +FOVY 35 +ITERATIONS 50 +FILE test.bmp +frame 0 +EYE 0 4.5 12 +VIEW 0 0 -1 +UP 0 1 0 +APERTURE 0 +FOCUSPLANE 10 + +OBJECT 0 +cube +material 0 +frame 0 +TRANS 0 0 0 +ROTAT 0 0 90 +SCALE .01 10 10 + +OBJECT 1 +cube +material 0 +frame 0 +TRANS 0 5 -5 +ROTAT 0 90 0 +SCALE .01 10 10 + +OBJECT 2 +cube +material 0 +frame 0 +TRANS 0 10 0 +ROTAT 0 0 90 +SCALE .01 10 10 + +OBJECT 3 +cube +material 1 +frame 0 +TRANS -5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +OBJECT 4 +cube +material 2 +frame 0 +TRANS 5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +OBJECT 5 +sphere +material 4 +frame 0 +TRANS 0 2 0 +ROTAT 0 180 0 +SCALE 3 3 3 + +OBJECT 6 +sphere +material 3 +frame 0 +TRANS 2 5 2 +ROTAT 0 180 0 +SCALE 2.5 2.5 2.5 + +OBJECT 7 +sphere +material 6 +frame 0 +TRANS -2 5 -2 +ROTAT 0 180 0 +SCALE 3 3 3 + + +OBJECT 8 +cube +material 8 +frame 0 +TRANS 0 10 0 +ROTAT 0 0 90 +SCALE .3 3 3 + +OBJECT 9 +sphere +material 7 +frame 0 +TRANS 3 8 0 +ROTAT 0 0 0 +SCALE 0.1 0.1 0.1 + +OBJECT 10 +cube +material 5 +frame 0 +TRANS 0 5 4 +ROTAT 0 5 0 +SCALE 2 2 2 diff --git a/scenes/sampleScene_simple.txt b/scenes/sampleScene_simple.txt new file mode 100644 index 0000000..0a6698f --- /dev/null +++ b/scenes/sampleScene_simple.txt @@ -0,0 +1,183 @@ +MATERIAL 0 //white diffuse +RGB 1 1 1 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 0 +REFRIOR 0 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 0 + +MATERIAL 1 //red diffuse +RGB .63 .06 .04 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 0 +REFRIOR 0 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 0 + +MATERIAL 2 //green diffuse +RGB .15 .48 .09 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 0 +REFRIOR 0 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 0 + +MATERIAL 3 //red glossy +RGB .63 .06 .04 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0.2 +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 1.0 +REFR 0 +REFRIOR 2 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 0 + +MATERIAL 5 //glass +RGB 1 1 1 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 1 +REFRIOR 2.2 +SCATTER 0 +ABSCOEFF .02 5.1 5.7 +RSCTCOEFF 13 +EMITTANCE 0 + +MATERIAL 6 //green glossy +RGB .15 .48 .09 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0.1 +REFR 0 +REFRIOR 2.6 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 0 + +MATERIAL 7 //light +RGB 1 0 0 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 0.8 + +MATERIAL 8 //light +RGB 0.25 0.65 0.85 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 15 + +CAMERA +RES 800 800 +FOVY 35 +ITERATIONS 500 +FILE fres.png +frame 0 +EYE 0 5 12 +VIEW 0 0 -1 +UP 0 1 0 +APERTURE 0 +FOCUSPLANE 10 + +OBJECT 0 +cube +material 0 +frame 0 +TRANS 0 0 0 +ROTAT 0 0 90 +SCALE .01 10 10 + +OBJECT 1 +cube +material 0 +frame 0 +TRANS 0 5 -5 +ROTAT 0 90 0 +SCALE .01 10 10 + +OBJECT 2 +cube +material 0 +frame 0 +TRANS 0 10 0 +ROTAT 0 0 90 +SCALE .01 10 10 + +OBJECT 3 +cube +material 1 +frame 0 +TRANS -5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +OBJECT 4 +cube +material 2 +frame 0 +TRANS 5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +OBJECT 5 +cube +material 8 +frame 0 +TRANS 0 5 0 +ROTAT 0 0 0 +SCALE 1 1 1 + +OBJECT 6 +sphere +material 5 +frame 0 +TRANS 0 2 -1 +ROTAT 0 0 0 +SCALE 2 2 2 + +OBJECT 7 +sphere +material 7 +frame 0 +TRANS 3 2 0 +ROTAT 0 0 0 +SCALE 2 2 2 \ No newline at end of file diff --git a/scenes/test1.txt b/scenes/test1.txt new file mode 100644 index 0000000..72f8026 --- /dev/null +++ b/scenes/test1.txt @@ -0,0 +1,175 @@ +MATERIAL 0 //white diffuse +RGB 1 1 1 +SPECEX 2 +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 2 +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 2 +SPECRGB 1 1 1 +REFL 0 +REFR 0 +REFRIOR 0 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 0 + +MATERIAL 3 //red glossy +RGB .63 .06 .04 +SPECEX 8 +SPECRGB 1 1 1 +REFL 1 +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 1 +REFR 0 +REFRIOR 2 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 0 + +MATERIAL 5 //glass +RGB .7 .6 1 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 1 +REFRIOR 5 +SCATTER 0 +ABSCOEFF .02 5.1 5.7 +RSCTCOEFF 13 +EMITTANCE 0 + +MATERIAL 6 //green glossy +RGB .15 .48 .09 +SPECEX 8 +SPECRGB 1 1 1 +REFL 1 +REFR 0 +REFRIOR 2.6 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 0 + +MATERIAL 7 //light +RGB 1 1 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 1 + +MATERIAL 8 //light +RGB 1 1 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 15 + +CAMERA +RES 800 800 +FOVY 25 +ITERATIONS 2500 +FILE test.bmp +frame 0 +EYE 0 4.5 12 +VIEW 0 0 -1 +UP 0 1 0 +APERTURE 0 +FOCUSPLANE 10 + +OBJECT 0 +cube +material 0 +frame 0 +TRANS 0 0 0 +ROTAT 0 0 90 +SCALE .01 10 10 + +OBJECT 1 +cube +material 0 +frame 0 +TRANS 0 5 -5 +ROTAT 0 90 0 +SCALE .01 10 10 + +OBJECT 2 +cube +material 0 +frame 0 +TRANS 0 10 0 +ROTAT 0 0 90 +SCALE .01 10 10 + +OBJECT 3 +cube +material 2 +frame 0 +TRANS -5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +OBJECT 4 +cube +material 1 +frame 0 +TRANS 5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +OBJECT 5 +cube +material 8 +frame 0 +TRANS 0 9.5 0 +ROTAT 0 0 90 +SCALE .3 3 3 + +OBJECT 6 +sphere +material 5 +frame 0 +TRANS 0 2 0 +ROTAT 0 180 0 +SCALE 3 3 3 \ No newline at end of file diff --git a/src/image.cpp b/src/image.cpp index 67bf157..46b6235 100755 --- a/src/image.cpp +++ b/src/image.cpp @@ -38,7 +38,7 @@ image::~image(){ //------------------------ float image::applyGamma(float f){ - //apply gamma correction, use simple power law gamma for now. + //apply gamma correction, use simple power law gamma for now. TODO: sRGB return pow(f/float(gamma.divisor), gamma.gamma); } diff --git a/src/interactions.h b/src/interactions.h index 6561796..98bf5aa 100755 --- a/src/interactions.h +++ b/src/interactions.h @@ -8,10 +8,56 @@ #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) { @@ -44,7 +90,81 @@ __host__ __device__ glm::vec3 calculateRandomDirectionInHemisphere(glm::vec3 nor //Now that you know how cosine weighted direction generation works, try implementing non-cosine (uniform) weighted random direction generation. //This should be much easier than if you had to implement calculateRandomDirectionInHemisphere. __host__ __device__ glm::vec3 getRandomDirectionInSphere(float xi1, float xi2) { - return glm::vec3(0,0,0); + + float theta = 2.0f * PI * xi1; + float phi = acos(2.0f * xi2 - 1); + + float radius = 1.0f; + + glm::vec3 localPosition = glm::vec3( radius * sin(theta) * cos(phi), radius * sin(theta) * sin(phi), radius * cos(phi)); + + return localPosition; + // No need to normalize since radius is one + // return glm::normalize(localPosition); } +//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){ + + // Reflected + if(fabs(m.hasReflective) > FLOAT_EPSILON) + { + ray reflected; + // reflected.direction = glm::reflect(r.direction,normal); + reflected.direction = r.direction - 2.0f * normal * glm::dot(normal,r.direction); + reflected.origin = intersect + NUDGE * reflected.direction; + r = reflected; + return 1; + } + // Transmitted + else if(fabs(m.hasRefractive) > FLOAT_EPSILON) + { + ray refracted; + + float dotValue = glm::dot(r.direction,normal); + float IOR = 1.0f/m.indexOfRefraction; + if(dotValue > 0) + { + IOR = 1.0f/IOR; + normal = -normal; + } + float cosValue = -glm::dot(r.direction,normal); + float k = 1 - IOR*IOR * (1 - cosValue * cosValue); + + // Total Internal Reflection + if(k < 0) + { + refracted.direction = r.direction - 2.0f * normal * glm::dot(normal,r.direction); + refracted.origin = intersect + NUDGE * refracted.direction; + r = refracted; + return 1; + } + + /* + float dotValue = glm::dot(r.direction,normal); + float IOR = 1.0f/m.indexOfRefraction; + if(dotValue > 0) + { + IOR = 1.0f/IOR; + normal = -normal; + } + + glm::vec3 refractedDirection = glm::refract(r.direction,normal,IOR); + refracted.direction = refractedDirection; + */ + + refracted.direction = glm::normalize(r.direction * IOR + normal * (IOR * cosValue - sqrt(k))); + refracted.origin = intersect + 0.001f * refracted.direction; + r = refracted; + return 2; + } + else // Diffuse + { + return 0; + } +}; + #endif diff --git a/src/intersections.h b/src/intersections.h index a6b9469..2515452 100755 --- a/src/intersections.h +++ b/src/intersections.h @@ -18,7 +18,6 @@ __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); @@ -69,98 +68,187 @@ __host__ __device__ glm::vec3 getSignOfRay(ray r){ return glm::vec3((int)(inv_direction.x < 0), (int)(inv_direction.y < 0), (int)(inv_direction.z < 0)); } -//Wrapper for cube intersection test for testing against unit cubes -__host__ __device__ float boxIntersectionTest(staticGeom box, ray r, glm::vec3& intersectionPoint, glm::vec3& normal){ - return boxIntersectionTest(glm::vec3(-.5,-.5,-.5), glm::vec3(.5,.5,.5), box, r, intersectionPoint, normal); -} - +//TODO: IMPLEMENT THIS FUNCTION //Cube intersection test, return -1 if no intersection, otherwise, distance to intersection -__host__ __device__ float boxIntersectionTest(glm::vec3 boxMin, glm::vec3 boxMax, staticGeom box, ray r, glm::vec3& intersectionPoint, glm::vec3& normal){ - 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); +__host__ __device__ float boxIntersectionTest(staticGeom box, ray r, glm::vec3& intersectionPoint, glm::vec3& normal){ + + glm::vec3 ro = multiplyMV(box.inverseTransform,glm::vec4(r.origin,1.0f)); + glm::vec3 rd = glm::normalize( multiplyMV(box.inverseTransform, glm::vec4(r.direction,0.0f)) ); + + ray rt; rt.origin = ro; rt.direction = rd; + + float halfWidth = 0.5f; + + + // Now box is in its own local unit space + glm::vec3 inverseDirection; + + glm::vec3 tmin, tmax; + + inverseDirection = 1.0f / rd; + + // min max x + if(inverseDirection.x >= 0) + { + tmin.x = (-halfWidth - rt.origin.x) * inverseDirection.x; + tmax.x = ( halfWidth - rt.origin.x) * inverseDirection.x; + } + else + { + tmin.x = ( halfWidth - rt.origin.x) * inverseDirection.x; + tmax.x = (-halfWidth - rt.origin.x) * inverseDirection.x; + } + + + // min max y + if(inverseDirection.y >= 0) + { + tmin.y = (-halfWidth - rt.origin.y) * inverseDirection.y; + tmax.y = ( halfWidth - rt.origin.y) * inverseDirection.y; + } + else + { + tmin.y = ( halfWidth - rt.origin.y) * inverseDirection.y; + tmax.y = (-halfWidth - rt.origin.y) * inverseDirection.y; + } + + if( (tmin.x > tmax.y) || (tmin.y > tmax.x)) return -1; + if( tmin.y > tmin.x ) tmin.x = tmin.y; + if( tmax.y < tmax.x ) tmax.x = tmax.y; + + // min max z + if(inverseDirection.z >= 0) + { + tmin.z = (-halfWidth - rt.origin.z) * inverseDirection.z; + tmax.z = ( halfWidth - rt.origin.z) * inverseDirection.z; + } + else + { + tmin.z = ( halfWidth - rt.origin.z) * inverseDirection.z; + tmax.z = (-halfWidth - rt.origin.z) * inverseDirection.z; + } + + if( (tmin.x > tmax.z) || (tmin.z > tmax.x)) return -1; + if( tmin.z > tmin.x ) tmin.x = tmin.z; + if( tmax.z < tmax.x ) tmax.x = tmax.z; + + float distanceLocal = glm::min(tmin.x,tmax.x); + + glm::vec3 localIntersectionPoint = getPointOnRay(rt, distanceLocal); + + glm::vec3 realIntersectionPoint = multiplyMV(box.transform, glm::vec4(localIntersectionPoint, 1.0)); + + intersectionPoint = realIntersectionPoint; + glm::vec3 localNormal; + + // Calculating Normal + if(fabs(localIntersectionPoint.x - -halfWidth) < FLOAT_EPSILON) + { + localNormal = glm::vec3(-1,0,0); + } + else if( fabs(localIntersectionPoint.x - halfWidth) < FLOAT_EPSILON) + { + localNormal = glm::vec3(1,0,0); + } + else if(fabs(localIntersectionPoint.y - -halfWidth) < FLOAT_EPSILON) + { + localNormal = glm::vec3(0,-1,0); + } + else if( fabs(localIntersectionPoint.y - halfWidth) < FLOAT_EPSILON) + { + localNormal = glm::vec3(0,1,0); + } + else if(fabs(localIntersectionPoint.z - -halfWidth) < FLOAT_EPSILON) + { + localNormal = glm::vec3(0,0,-1); + } + else if( fabs(localIntersectionPoint.z - halfWidth) < FLOAT_EPSILON) + { + localNormal = glm::vec3(0,0,1); + } + + + // Psuedo point is one local unit distance behind the local intersection point in the direction of the normal + glm::vec3 realPseudoPoint = multiplyMV(box.transform, glm::vec4(localIntersectionPoint - localNormal ,1.0f)); + normal = glm::normalize(realIntersectionPoint - realPseudoPoint); + + // @DO: something is not normalized. Confirm what + //return glm::length(r.origin - realIntersectionPoint); + //return glm::length(r.origin - realIntersectionPoint) * glm::length(r.direction); + + return distanceLocal; + + + + /* + 560- ray tracer implementation + + float tNear = -1000000, tFar = 10000000; + float t1,t2; + + float returnT; + + for(int i=0; i<3;i++) + { + float xD, xo; + switch(i) + { + case 0: xD = rt.direction.x; xo = rt.origin.x; break; + case 1: xD = rt.direction.y; xo = rt.origin.y; break; + case 2: xD = rt.direction.z; xo = rt.origin.z; break; + } + if(fabs(xD)<0.00001) + { + if(xo < -halfWidth || xo > halfWidth) + return -1; // Handle Parallel Case + } + t1 = (-halfWidth - xo)/xD; //(-0.5 - xo)/xD; // my implementation has -1 to 1 + t2 = ( halfWidth - xo)/xD; //( 0.5 - xo)/xD; // my implementation has -1 to 1 + + + //@DO: is this how insides are checked? + if ( t1 > t2)// && !(t2<0)) + { + float temp = t2; + t2 = t1; + t1 = temp; + } + if (t1 > tNear ) + tNear = t1; + if (t2 < tFar)// && t2>0) + tFar = t2; + if( tNear > tFar) + return -1;//Box missed, do something + if( tFar < 0 ) + return -1;//Box behind ray, do something + + if(tNear < 0) + returnT = tFar; + else + returnT = tNear; + } + //cout< u01(-1,1); - thrust::uniform_real_distribution u02(0,TWO_PI); - - float theta = (float)u02(rng); - float cosphi = (float)u01(rng); - float sinphi = sqrt(1 - cosphi*cosphi); - glm::vec3 point = radius*glm::vec3(sinphi*cos(theta),sinphi*sin(theta),cosphi); - glm::vec3 randPoint = multiplyMV(sphere.transform, glm::vec4(point,1.0f)); - - return randPoint; + + /* + thrust::default_random_engine rng(hash(randomSeed)); + thrust::uniform_real_distribution uniformDistribution(-1,1); + float x1, x2; + int flag = 0; + while(flag == 0) + { + x1 = (float) uniformDistribution(rng); + x2 = (float) uniformDistribution(rng); + + if( (x1*x1 + x2*x2) < 1.0f) + flag = 1; + } + + float term1 = x1*x1 + x2*x2; + float term2 = sqrt(1 - term1); + + glm::vec3 unitSphere = glm::vec3( 2 * x1 * term2, + 2 * x2 * term2, + 1 - 2 * term1 ); + + */ + + thrust::default_random_engine rng(hash(randomSeed)); + thrust::uniform_real_distribution uniformDistribution(0,1); + float x1, x2; + x1 = (float) uniformDistribution(rng); + x2 = (float) uniformDistribution(rng); + + float theta = 2.0f * PI * x1; + float phi = acos(2.0f * x2 - 1); + + float radius = 0.5f; + + glm::vec3 localPosition = glm::vec3( radius * cos(theta) * sin(phi), radius * sin(theta) * sin(phi), radius * cos(phi)); + + glm::vec3 realPosition = multiplyMV(sphere.transform, glm::vec4(localPosition,1.0f)); + + return realPosition; } #endif diff --git a/src/main.cpp b/src/main.cpp index 81836b1..1be0bd8 100755 --- a/src/main.cpp +++ b/src/main.cpp @@ -13,6 +13,9 @@ int main(int argc, char** argv){ + liveCamera.aperture = 0;//1.75; + liveCamera.focusPlane = 0;//12.5; + #ifdef __APPLE__ // Needed in OSX to force use of OpenGL3.2 glfwOpenWindowHint(GLFW_OPENGL_VERSION_MAJOR, 3); @@ -26,7 +29,7 @@ int main(int argc, char** argv){ finishedRender = false; targetFrame = 0; - singleFrameMode = false; + singleFrameMode = true; // Load scene file for(int i=1; iiterations){ uchar4 *dptr=NULL; iterations++; @@ -123,7 +142,7 @@ void runCuda(){ // execute the kernel - cudaRaytraceCore(dptr, renderCam, targetFrame, iterations, materials, renderScene->materials.size(), geoms, renderScene->objects.size() ); + cudaRaytraceCore(dptr, renderCam, targetFrame, iterations, materials, renderScene->materials.size(), geoms, renderScene->objects.size(), liveCamera ); // unmap buffer object cudaGLUnmapBufferObject(pbo); @@ -142,8 +161,8 @@ void runCuda(){ gammaSettings gamma; gamma.applyGamma = true; - gamma.gamma = 1.0; - gamma.divisor = 1.0; //renderCam->iterations; + gamma.gamma = 1.0;//2.2; + gamma.divisor = renderCam->iterations; outputImage.setGammaSettings(gamma); string filename = renderCam->imageName; string s; @@ -172,7 +191,13 @@ void runCuda(){ finishedRender = false; } } - + // Performance Analysis End + cudaEventRecord(stop, 0); // Trigger Stop event + cudaEventSynchronize(stop); // Sync events (BLOCKS till last (stop in this case) has been recorded!) + float elapsedTime; // Initialize elapsedTime; + cudaEventElapsedTime(&elapsedTime, start, stop); // Calculate runtime, write to elapsedTime -- cudaEventElapsedTime returns value in milliseconds. Resolution ~0.5ms + + printf("Execution Time: %fms for iteration: %d\n", elapsedTime, iterations); // Print Elapsed time } #ifdef __APPLE__ @@ -199,9 +224,10 @@ void runCuda(){ #else void display(){ + runCuda(); - string title = "565Raytracer | " + utilityCore::convertIntToString(iterations) + " Iterations"; + string title = "565PathTracer | " + utilityCore::convertIntToString(iterations) + " Iterations"; glutSetWindowTitle(title.c_str()); glBindBuffer( GL_PIXEL_UNPACK_BUFFER, pbo); @@ -218,17 +244,101 @@ void runCuda(){ glutSwapBuffers(); } + void resetAccumulator() + { + iterations = 1; + } + void keyboard(unsigned char key, int x, int y) { - std::cout << key << std::endl; +// std::cout << key << std::endl; switch (key) { case(27): exit(1); break; + case('w'): + iterations = 1; + cudaClearAccumulatorImage(renderCam); + liveCamera.position.y += STRAFE_AMOUNT; + break; + case('s'): + iterations = 1; + cudaClearAccumulatorImage(renderCam); + liveCamera.position.y -= STRAFE_AMOUNT; + break; + case('a'): + iterations = 1; + cudaClearAccumulatorImage(renderCam); + liveCamera.position.x += STRAFE_AMOUNT; + break; + case('d'): + iterations = 1; + cudaClearAccumulatorImage(renderCam); + liveCamera.position.x -= STRAFE_AMOUNT; + break; + case('q'): + iterations = 1; + cudaClearAccumulatorImage(renderCam); + liveCamera.position.z += STRAFE_AMOUNT; + break; + case('z'): + iterations = 1; + cudaClearAccumulatorImage(renderCam); + liveCamera.position.z -= STRAFE_AMOUNT; + break; + case('['): + iterations = 1; + cudaClearAccumulatorImage(renderCam); + liveCamera.view.x += STRAFE_AMOUNT; + break; + case(']'): + iterations = 1; + cudaClearAccumulatorImage(renderCam); + liveCamera.view.x -= STRAFE_AMOUNT; + break; + case('o'): + iterations = 1; + cudaClearAccumulatorImage(renderCam); + liveCamera.view.y += STRAFE_AMOUNT; + break; + case('p'): + iterations = 1; + cudaClearAccumulatorImage(renderCam); + liveCamera.view.y -= STRAFE_AMOUNT; + break; } } + void specialKeyboard(int key, int x, int y) + { + switch (key) + { + case GLUT_KEY_UP: + iterations = 1; + cudaClearAccumulatorImage(renderCam); + liveCamera.focusPlane += STRAFE_AMOUNT; + break; + case GLUT_KEY_DOWN: + iterations = 1; + cudaClearAccumulatorImage(renderCam); + liveCamera.focusPlane -= STRAFE_AMOUNT; + break; + case GLUT_KEY_LEFT: + iterations = 1; + cudaClearAccumulatorImage(renderCam); + liveCamera.aperture += 0.5*STRAFE_AMOUNT; + break; + case GLUT_KEY_RIGHT: + iterations = 1; + cudaClearAccumulatorImage(renderCam); + liveCamera.aperture -= 0.5*STRAFE_AMOUNT; + if(liveCamera.aperture < NUDGE) + liveCamera.aperture = NUDGE; + break; + } + } + #endif diff --git a/src/main.h b/src/main.h index 0bab7cb..0327bed 100755 --- a/src/main.h +++ b/src/main.h @@ -62,6 +62,7 @@ GLuint texcoordsLocation = 1; const char *attributeLocations[] = { "Position", "Tex" }; GLuint pbo = (GLuint)NULL; GLuint displayImage; +cameraData liveCamera; //------------------------------- //----------CUDA STUFF----------- @@ -86,6 +87,7 @@ void runCuda(); #else void display(); void keyboard(unsigned char key, int x, int y); + void specialKeyboard(int key, int x, int y); #endif //------------------------------- diff --git a/src/raytraceKernel.cu b/src/raytraceKernel.cu index 87a65a6..cfd15ef 100755 --- a/src/raytraceKernel.cu +++ b/src/raytraceKernel.cu @@ -9,17 +9,37 @@ #include #include #include "sceneStructs.h" +#include "glm/glm.hpp" #include "utilities.h" #include "raytraceKernel.h" #include "intersections.h" #include "interactions.h" #include -#include "glm/glm.hpp" + +// Optimizations and add ons +#define JITTER 1 +#define COMPACTION 1 +#define ACCUMULATION 1 +#define DOF 1 +#define FRESNEL 1 +#define SCHLICK 1 +#define PATHTRACER 1 +#define PAINTERLY 0 + +#if CUDA_VERSION >= 5000 + #include +#else + #include +#endif + +glm::vec3* accumulatorImage = NULL; +extern bool singleFrameMode; void checkCUDAError(const char *msg) { cudaError_t err = cudaGetLastError(); if( cudaSuccess != err) { fprintf(stderr, "Cuda error: %s: %s.\n", msg, cudaGetErrorString( err) ); + getchar(); exit(EXIT_FAILURE); } } @@ -35,41 +55,52 @@ __host__ __device__ glm::vec3 generateRandomNumberFromThread(glm::vec2 resolutio return glm::vec3((float) u01(rng), (float) u01(rng), (float) u01(rng)); } -//Kernel that does the initial raycast from the camera. -__host__ __device__ ray raycastFromCameraKernel(glm::vec2 resolution, float time, int x, int y, glm::vec3 eye, glm::vec3 view, glm::vec3 up, glm::vec2 fov){ - - 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! +//TODO: IMPLEMENT THIS FUNCTION +//Function that does the initial raycast from the camera +__host__ __device__ ray raycastFromCameraKernel(glm::vec2 resolution, float time, int x, int y, glm::vec3 eye, glm::vec3 view, glm::vec3 up, glm::vec2 fov +#if DOF + ,float aperture, float focusPlane +#endif + ){ ray r; + + // @DO: verify field of view! + glm::vec3 axis_a = glm::cross(view, up); + glm::vec3 axis_b = glm::cross(axis_a, view); + glm::vec3 midPoint = eye + view; + glm::vec3 viewPlaneX = axis_a * tan(PI_F * fov.x/180.0f) * glm::length(view)/glm::length(axis_a); + glm::vec3 viewPlaneY = axis_b * tan(PI_F * fov.y/180.0f) * glm::length(view)/glm::length(axis_b); + +#if JITTER + glm::vec3 jitter = generateRandomNumberFromThread(resolution,time,x,y); + glm::vec3 screenPoint = midPoint + + (2.0f * ((jitter.x + 1.0f * x) / (resolution.x-1)) - 1.0f) * viewPlaneX + + (1.0f - 2.0f * ((jitter.y + 1.0f * y) / (resolution.y-1))) * viewPlaneY; +#else + glm::vec3 screenPoint = midPoint + + (2.0f * (1.0f * x / (resolution.x-1)) - 1.0f) * viewPlaneX + + (1.0f - 2.0f * (1.0f * y / (resolution.y-1))) * viewPlaneY; + +#endif + +#if DOF + + glm::vec3 focusPlaneIntersection; + r.origin = eye; - r.direction = direction; + r.direction = glm::normalize(screenPoint - eye); + + glm::vec3 focusPlaneCenter = r.origin + r.direction * focusPlane; + planeIntersectionTest(focusPlaneCenter,view,r,focusPlaneIntersection); + + glm::vec3 apertureJitter = aperture * (generateRandomNumberFromThread(resolution,time,x,y) - 0.5f); + r.origin = r.origin + apertureJitter; + r.direction = glm::normalize(focusPlaneIntersection - r.origin); + +#else + r.origin = screenPoint; + r.direction = glm::normalize(screenPoint - eye); +#endif return r; } @@ -83,8 +114,8 @@ __global__ void clearImage(glm::vec2 resolution, glm::vec3* image){ } } -//Kernel that writes the image to the OpenGL PBO directly. -__global__ void sendImageToPBO(uchar4* PBOpos, glm::vec2 resolution, glm::vec3* image){ +//Kernel that writes the image to the OpenGL PBO directly. +__global__ void sendImageToPBO(uchar4* PBOpos, glm::vec2 resolution, glm::vec3* image, float frames){ int x = (blockIdx.x * blockDim.x) + threadIdx.x; int y = (blockIdx.y * blockDim.y) + threadIdx.y; @@ -92,10 +123,10 @@ __global__ void sendImageToPBO(uchar4* PBOpos, glm::vec2 resolution, glm::vec3* 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; + glm::vec3 color; + color.x = image[index].x*255.0 / frames; + color.y = image[index].y*255.0 / frames; + color.z = image[index].z*255.0 / frames; if(color.x>255){ color.x = 255; @@ -111,64 +142,668 @@ __global__ void sendImageToPBO(uchar4* PBOpos, glm::vec2 resolution, glm::vec3* // Each thread writes one pixel location in the texture (textel) PBOpos[index].w = 0; - PBOpos[index].x = color.x; + PBOpos[index].x = color.x; PBOpos[index].y = color.y; PBOpos[index].z = color.z; } } +__device__ bool visibilityCheck(ray r, staticGeom* geoms, int numberOfGeoms, glm::vec3 pointToCheck, int geomShotFrom, int lightSourceIndex) +{ + bool visible = true; + float distance = glm::length(r.origin - pointToCheck); + + // Check whether any object occludes point to check from ray's origin + for(int iter=0; iter < numberOfGeoms; iter++) + { + // Avoid calculating self intersections + if(iter==lightSourceIndex) + continue; + + float depth=-1; + glm::vec3 intersection; + glm::vec3 normal; + + if(geoms[iter].type == CUBE) + { + depth = boxIntersectionTest(geoms[iter],r,intersection,normal); + } + + + else if(geoms[iter].type == SPHERE) + { + depth = sphereIntersectionTest(geoms[iter],r,intersection,normal); + } + + if(depth > 0 && (depth + NUDGE) < distance) + { + //printf("Depth: %f\n", depth); + visible = false; + break; + } + } + + + return visible; +} + //TODO: IMPLEMENT THIS FUNCTION //Core raytracer kernel -__global__ void raytraceRay(glm::vec2 resolution, float time, float bounce, cameraData cam, int rayDepth, glm::vec3* colors, - staticGeom* geoms, int numberOfGeoms, material* materials, int numberOfMaterials){ +__global__ void raytraceRay(glm::vec2 resolution, float time, cameraData cam, int rayDepth, glm::vec3* colors, + staticGeom* geoms, int numberOfGeoms, material* materials, int numberOfMaterials, ray* rayPool +#if COMPACTION + , int numberOfRays +#endif + ){ + int index; + int x,y; +#if COMPACTION + index = blockIdx.x * blockDim.x + threadIdx.x; + y = index / resolution.x; + x = index - y * resolution.x; + if (index < numberOfRays) { +#else + x = (blockIdx.x * blockDim.x) + threadIdx.x; + y = (blockIdx.y * blockDim.y) + threadIdx.y; + index = x + (y * resolution.x); + if((x<=resolution.x && y<=resolution.y) && glm::length(rayPool[index].transmission) > FLOAT_EPSILON){ +#endif + + ray r = rayPool[index]; + + + //Check all geometry for intersection + int intersectedGeom = -1; + int intersectedMaterial = -1; + float minDepth = 1000000.0f; + glm::vec3 minIntersectionPoint; + glm::vec3 minNormal = glm::vec3(0.0f); + for(int iter=0; iter < numberOfGeoms; iter++) + { + float depth=-1; + glm::vec3 intersection; + glm::vec3 normal; + staticGeom currentGeometry = geoms[iter]; + if(currentGeometry.type == CUBE) + { + depth = boxIntersectionTest(currentGeometry,r,intersection,normal); + } + + else if(geoms[iter].type == SPHERE) + { + depth = sphereIntersectionTest(currentGeometry,r,intersection,normal); + } + + + if(depth > 0 && depth < minDepth) + { + minDepth = depth; + minIntersectionPoint = intersection; + minNormal = normal; + intersectedGeom = iter; + intersectedMaterial = currentGeometry.materialid; + } + } + + // Depth render - test + //float maxDepth = 15.0f; + + glm::vec3 diffuseLight = glm::vec3(0.0f); + glm::vec3 phongLight = glm::vec3(0.0f); + + glm::vec3 diffuseColor; + glm::vec3 specularColor; + glm::vec3 emittance; + + //Calculate Lighting if any geometry is intersected + if(intersectedGeom > -1) + { + //finalColor = materials[geoms[intersectedGeom].materialid].color; + material m = materials[intersectedMaterial]; + diffuseColor = m.color; + specularColor = m.specularColor; + // Emmited color is the same as material color + emittance = m.color * m.emittance; + +#if PATHTRACER != 1 + // Stochastic Diffused Lighting with "area" lights + for(int iter = 0; iter < numberOfGeoms; iter++) + { + material lightMaterial = materials[geoms[iter].materialid]; + // If this geometry is going to act like a light source + if(lightMaterial.emittance > 0.0001f) + { + glm::vec3 lightSourceSample; + + // Get a random point on the light source + if(geoms[iter].type == SPHERE) + { + lightSourceSample = getRandomPointOnSphere(geoms[iter],time*index); + } + else if(geoms[iter].type == CUBE) + { + lightSourceSample = getRandomPointOnCube(geoms[iter],time*index); + } + + // Diffuse Lighting Calculation + glm::vec3 L = glm::normalize(lightSourceSample - minIntersectionPoint); + + //Shadow Ray check + ray shadowRay; + shadowRay.origin = minIntersectionPoint + NUDGE * L; + shadowRay.direction = L; + + bool visible = visibilityCheck(shadowRay,geoms,numberOfGeoms,lightSourceSample, intersectedGeom, iter); + + if(visible) + { + diffuseLight += lightMaterial.color * lightMaterial.emittance * glm::max(glm::dot(L,minNormal),0.0f); + + + // Calculate Phong Specular Part only if exponent is greater than 0 + if(m.specularExponent > FLOAT_EPSILON) + { + glm::vec3 reflectedLight = 2.0f * minNormal * glm::dot(minNormal, L) - L; + phongLight += lightMaterial.color * lightMaterial.emittance * pow(glm::max(glm::dot(reflectedLight,minNormal),0.0f),m.specularExponent); + } + + } + } + } +#endif; + + AbsorptionAndScatteringProperties absScatProps; + glm::vec3 colorSend, unabsorbedColor; + ray returnRay = r; + int rayPropogation = calculateBSDF(returnRay,minIntersectionPoint,minNormal,diffuseColor*m.emittance,absScatProps,colorSend,unabsorbedColor,m); + + // Diffuse Reflection or light source + if (rayPropogation == 0) + { +#if PATHTRACER +#if COMPACTION + colors[r.pixelIndex] += r.transmission * emittance; +#else + colors[index] += r.transmission * emittance; +#endif +#if PAINTERLY + glm::vec3 randomVector = generateRandomNumberFromThread(resolution,time * (rayDepth+1),1,1); +#else + glm::vec3 randomVector = generateRandomNumberFromThread(resolution,time * (rayDepth+1),x,y); +#endif + r.direction = calculateRandomDirectionInHemisphere(minNormal, randomVector.x, randomVector.y); + r.origin = minIntersectionPoint + 0.0005f * r.direction; + r.transmission *= diffuseColor; + +#else +#if COMPACTION + colors[r.pixelIndex] += r.transmission * ( emittance + diffuseLight * diffuseColor + phongLight * specularColor); +#else + colors[index] += r.transmission * ( emittance + diffuseLight * diffuseColor + phongLight * specularColor); +#endif + r.transmission = glm::vec3(0); +#endif + rayPool[index] = r; + + } + // Reflection; calculate transmission coeffiecient + else if(rayPropogation == 1) + { +#if PATHTRACER +#if COMPACTION + colors[r.pixelIndex] += r.transmission * emittance; +#else + colors[index] += r.transmission * emittance; +#endif +#else +#if COMPACTION + colors[r.pixelIndex] += r.transmission * (1.0f - m.hasReflective) * ( emittance + diffuseLight * diffuseColor + phongLight * specularColor); +#else + colors[index] += r.transmission * (1.0f - m.hasReflective) * ( emittance + diffuseLight * diffuseColor + phongLight * specularColor); +#endif +#endif + returnRay.transmission = r.transmission * diffuseColor * m.hasReflective; + rayPool[index] = returnRay; + } + // Refraction; calculate transmission coeffiecient + else if (rayPropogation == 2) + { + +#if PATHTRACER +#if COMPACTION + colors[r.pixelIndex] += r.transmission * emittance; +#else + colors[index] += r.transmission * emittance; +#endif +#else +#if COMPACTION + colors[r.pixelIndex] += r.transmission * (1.0f - m.hasRefractive) * ( emittance + diffuseLight * diffuseColor + phongLight * specularColor); +#else + colors[index] += r.transmission * (1.0f - m.hasRefractive) * ( emittance + diffuseLight * diffuseColor + phongLight * specularColor); +#endif +#endif + returnRay.transmission = r.transmission * diffuseColor * m.hasRefractive; + + +#if FRESNEL + // Fresnel Calculation + + // Fabs because the angle is always between 0 and 90, direction not-withstanding + float nd = fabs(glm::dot(r.direction, minNormal)); + float nt = fabs(glm::dot(returnRay.direction, minNormal)); + float n_a = nd < 0 ? 1.0f : m.indexOfRefraction; + float n_b = nd < 0 ? m.indexOfRefraction : 1.0f; + float amountReflected; + +#if SCHLICK + // Schlick's Approximation + + float RO = (n_a - n_b) * (n_a - n_b) / ( (n_a + n_b) * (n_a + n_b)); + float c; + if(n_a < n_b) + c = 1 - nd; + else + c = 1 - nt; + + amountReflected = RO + (1-RO) * c * c * c * c * c; + +#else + // Fresnels equations + float reflectedParallel = (n_b * nd - n_a * nt) * (n_b * nd - n_a * nt) / ((n_b * nd + n_a * nt) * (n_b * nd + n_a * nt)); + float reflectedPerpendicular = (n_a * nd - n_b * nt) * (n_a * nd - n_b * nt) / ((n_a * nd + n_b * nt) * (n_a * nd + n_b * nt)); + amountReflected = 0.5 * (reflectedParallel + reflectedPerpendicular); +#endif + // Stochastically decide whether to reflect or refract + glm::vec3 randVector = generateRandomNumberFromThread(resolution,time * (rayDepth+1),x,y); + + // If a uniform variable is less than the reflected amount, this ray shall be reflected + if(randVector.y < amountReflected) + { + returnRay.direction = r.direction - 2.0f * minNormal * glm::dot(minNormal,r.direction); + returnRay.origin = minIntersectionPoint + NUDGE * returnRay.direction; + } +#endif + rayPool[index] = returnRay; + } + } + // No intersection, mark rays as dead + // Ambeint term + else + { + glm::vec3 ambient = glm::vec3(0,0,0); +#if COMPACTION + colors[r.pixelIndex] += ambient; +#else + colors[index] += ambient; +#endif + r.transmission = glm::vec3(0.0f); + rayPool[index] = r; + } + + /* + //Checking for correct ray direction + colors[index].x = fabs(r.direction.x); + colors[index].y = fabs(r.direction.y); + colors[index].z = fabs(r.direction.z); + + //Check for correct material pickup + colors[index] = color; + + //Checking for correct depth testing + colors[index] = color * (maxDepth - minDepth)/maxDepth; + + //Checking for correct normals + colors[index] = glm::vec3(minNormal); + colors[index] = glm::vec3( fabs(minNormal.x), fabs(minNormal.y), fabs(minNormal.z)); + */ + + + + } +} + +__global__ void fillRayPoolFromCamera(glm::vec2 resolution, float time, cameraData cam, int rayDepth, glm::vec3* colors, + staticGeom* geoms, int numberOfGeoms, ray* rayPool){ 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)){ + + + ray r; + r = raycastFromCameraKernel(resolution,time,x,y,cam.position,cam.view,cam.up,cam.fov +#if DOF + ,cam.aperture, cam.focusPlane +#endif + ); + +#if COMPACTION + r.pixelIndex = index; +#endif + r.transmission = glm::vec3(1.0f); + + // Access global memory only once + rayPool[index] = r; + } +} + +__global__ void combineIntoAccumulatorImage(glm::vec2 resolution, float frames, glm::vec3* inputColors, glm::vec3* displayColors) +{ + 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)){ + //displayColors[index] = (((frames-1) * displayColors[index]) + inputColors[index])/frames; - float MAX_DEPTH = 100000000000000000; - float depth = MAX_DEPTH; - - for(int i=0; i-EPSILON){ - MAX_DEPTH = depth; - colors[index] = materials[geoms[i].materialid].color; - } - } + // Averaging done in SendToPBO + displayColors[index] += inputColors[index]; + } +} +#if COMPACTION - //colors[index] = generateRandomNumberFromThread(resolution, time, x, y); - } +/*** + + Compaction: + - Mark whether ray should be compacted or not based on transmission + - Inclusive Scan over predicate array (in place?) to generate output indices + Now all rays corresponding to an ouput index > 0 should go to index outputIndex - 1 in output array + - Use a scatter operation to distribute ray data onto those locations in the output array. + +***/ + +// NVidiaScan +__global__ void prescan(float *g_odata, float *g_idata, int n) +{ + + int bid = blockIdx.x; + + // must give final sum into auxilary array at bid. + // maximum number of elements that can be prescanned? + // number of threads per block * number of threads per block (each scan per block and then one block to scan all the in between things) + // @TODO: how to scan more than the above described limit? + + // NVIDIA implenentation follows from: + // http://http.developer.nvidia.com/GPUGems3/gpugems3_ch39.html + + extern __shared__ float temp[];// allocated on invocation + int thid = threadIdx.x; + int offset = 1; + temp[2*thid] = g_idata[2*thid]; // load input into shared memory + temp[2*thid+1] = g_idata[2*thid+1]; + for (int d = n>>1; d > 0; d >>= 1) // build sum in place up the tree + { + __syncthreads(); + if (thid < d) + { + int ai = offset*(2*thid+1)-1; + int bi = offset*(2*thid+2)-1; + temp[bi] += temp[ai]; + } + offset *= 2; + } + if (thid == 0) { temp[n - 1] = 0; } // clear the last element + for (int d = 1; d < n; d *= 2) // traverse down tree & build scan + { + offset >>= 1; + __syncthreads(); + if (thid < d) + { + int ai = offset*(2*thid+1)-1; + int bi = offset*(2*thid+2)-1; + float t = temp[ai]; + temp[ai] = temp[bi]; + temp[bi] += t; + } + } + __syncthreads(); + g_odata[2*thid] = temp[2*thid]; // write results to device memory + g_odata[2*thid+1] = temp[2*thid+1]; } +// Mark with predicate whether active or inactive +__global__ void predicateMark(ray* inputRays, int* outputPredicate, int size) +{ + // Using 1D kernel for compaction + int index = (blockIdx.x * blockDim.x) + threadIdx.x; + + if(index < size) + { + if(glm::length(inputRays[index].transmission) > FLOAT_EPSILON) + { + outputPredicate[index] = 1; + } + else + { + outputPredicate[index] = 0; + } + } +} + +// Scan Per Block +__global__ void naiveScanPerBlock(int *inData, int* outData, int *blockSum, int size) +{ + int index = blockIdx.x * blockDim.x + threadIdx.x; + int threadId = threadIdx.x; + + // third arg of kernel launch = numFloats*sizeofFloat + extern __shared__ int sharedElements[]; + + // Fill elements into shared Memory + if(index < size) + sharedElements[threadId] = inData[index]; + // If last block, fill in zeros for padding + else + sharedElements[threadId] = 0; + + __syncthreads(); + + int numSteps = ceil(log(1.0f*blockDim.x)/log(2.0f)); + int offset = 1; + for(int d = 0; d < numSteps; d++, offset *= 2) + { + int t = sharedElements[threadId]; + if(threadId >= offset) + { + t = t + sharedElements[threadId - offset]; + } + __syncthreads(); + + sharedElements[threadId] = t; + __syncthreads(); + } + + outData[index] = sharedElements[threadId]; + if(threadIdx.x == blockDim.x-1) + { + blockSum[blockIdx.x] = sharedElements[threadId]; + } +} + +// Add back Scanned blocksums +__global__ void addBackBlockSums(int *outData, int *blockSum, int size, int *returnedSum) +{ + int index = blockIdx.x * blockDim.x +threadIdx.x; + if(blockIdx.x > 0) + outData[index] = outData[index] + blockSum[blockIdx.x-1]; + if(index == size - 1) + returnedSum[0] = outData[index]; +} + +// Wrap around total scanned array +int parallelScan(int *inData, int *outData, int size, int d=0) +{ + if(size==0) + return 0; + + int threads = 1024; + int blocks = ceil(size*1.0f/threads); + dim3 blocksPerGrid(blocks,1,1); + dim3 threadsPerBlock(threads,1,1); + + //std::cout<<"ParallelScan in "<>>(inData, outData,cudaBlockSum,size); + checkCUDAError("Naive Scan Failed!"); + + // If the number of blocks is 1 + //if() + // only do naive scan per block for one block and copy into cudaBlockSumScan + // naiveScanPerBlock<<<1,threadsPerBlock, threads*sizeof(float)>>>(....) + //else recursive + // parallelScan(cudaBlockSum,cudaBlockSumScan,blocks); + + // Base Case of recursion + if (blocks==1) + { + cudaMemcpy(cudaSum,cudaBlockSum,sizeof(int),cudaMemcpyDeviceToDevice); + //return 0; + } + else + { + parallelScan(cudaBlockSum,cudaBlockSumScan,blocks,d+1); + addBackBlockSums<<>>(outData, cudaBlockSumScan, size, cudaSum); + checkCUDAError("Add Back Blocks Failed!"); + } + + + if(d==0) + { + cudaMemcpy(sum,cudaSum,sizeof(int),cudaMemcpyDeviceToHost); + } + + cudaFree(cudaSum); + cudaFree(cudaBlockSumScan); + cudaFree(cudaBlockSum); + return (*sum); + +} + +// Scatter Rays into the appropriate locations in the output array +__global__ void scatter(ray* inputRays, ray* outputRays, int* predicate, int* scatterIndices, int size) +{ + // Using 1D kernel for compaction + int index = (blockIdx.x * blockDim.x) + threadIdx.x; + + if(index < size) + { + int scatterIndex = scatterIndices[index]; + // Data with scatter index < 0 is data not to be used + if(predicate[index] > 0) + { + outputRays[scatterIndex - 1] = inputRays[index]; + } + } +} + +// Stream Compaction +int streamCompactRayPool(ray* inputRays, ray* outputRays, int size) +{ + + /* + // Testing recursive scan primitive + + int inputData[] = {1,1,1,1,1,1,1,1};//{1,1,2,3,4,5,6,7}; + int dataSize = 8; + + int *cudaInputData; + cudaMalloc((void**)&cudaInputData,dataSize*sizeof(int)); + cudaMemcpy(cudaInputData,inputData,dataSize*sizeof(int),cudaMemcpyHostToDevice); + + int *cudaOutputData; + cudaMalloc((void**)&cudaOutputData,dataSize*sizeof(int)); + + int val = parallelScan(cudaInputData,cudaOutputData,dataSize); + + + std::cout<<"Number of elements: "<>>(inputRays, predicateArray, size); + checkCUDAError("Predicate Mark Failed!"); + + // Scan Predicate to get location and also total number of final rays + int compactedSize = parallelScan(predicateArray,scatterLocations,size); + + // Scatter rays to new locations in output array + scatter<<>>(inputRays,outputRays,predicateArray,scatterLocations,size); + checkCUDAError("Scatter Failed!"); + + cudaFree(scatterLocations); + cudaFree(predicateArray); + + return compactedSize; +} + + +#endif + //TODO: FINISH THIS FUNCTION // Wrapper for the __global__ call that sets up the kernel calls and does a ton of memory management -void cudaRaytraceCore(uchar4* PBOpos, camera* renderCam, int frame, int iterations, material* materials, int numberOfMaterials, geom* geoms, int numberOfGeoms){ - +void cudaRaytraceCore(uchar4* PBOpos, camera* renderCam, int frame, int iterations, material* materials, int numberOfMaterials, geom* geoms, int numberOfGeoms, cameraData liveCamera){ + + + // testing + //streamCompact(); + + if(iterations == 1) + { + // Allocate Accumulator Image + cudaAllocateAccumulatorImage(renderCam); + } + 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)); @@ -192,36 +827,119 @@ void cudaRaytraceCore(uchar4* PBOpos, camera* renderCam, int frame, int iteratio cudaMalloc((void**)&cudageoms, numberOfGeoms*sizeof(staticGeom)); cudaMemcpy( cudageoms, geomList, numberOfGeoms*sizeof(staticGeom), cudaMemcpyHostToDevice); - material* cudamaterials = NULL; - cudaMalloc((void**)&cudamaterials, numberOfMaterials*sizeof(material)); - cudaMemcpy( cudamaterials, materials, numberOfMaterials*sizeof(material), cudaMemcpyHostToDevice); - //package camera cameraData cam; cam.resolution = renderCam->resolution; cam.position = renderCam->positions[frame]; cam.view = renderCam->views[frame]; cam.up = renderCam->ups[frame]; + cam.aperture = renderCam->apertures[frame]; + cam.focusPlane = renderCam->focusPlanes[frame]; cam.fov = renderCam->fov; - //kernel launches - for(int bounce = 1; bounce <= 1; ++bounce) + //user interaction + cam.position += (liveCamera.position); + cam.view = glm::normalize(cam.view + liveCamera.view); + cam.aperture += liveCamera.aperture; + cam.focusPlane += liveCamera.focusPlane; + + //Transfer materials + material* cudamaterials = NULL; + cudaMalloc((void**)&cudamaterials, numberOfMaterials*sizeof(material)); + cudaMemcpy( cudamaterials, materials, numberOfMaterials*sizeof(material), cudaMemcpyHostToDevice); + + //Create Memory for RayPool + ray* cudarays = NULL; + cudaMalloc((void**)&cudarays, (renderCam->resolution.x * renderCam->resolution.y) * sizeof(ray)); + +#if COMPACTION + ray* cudarays2 = NULL; + cudaMalloc((void**)&cudarays2, (renderCam->resolution.x * renderCam->resolution.y) * sizeof(ray)); +#endif + + //clear On screen buffer + clearImage<<>>(renderCam->resolution, cudaimage); + + //Fill ray pool with rays from camera for first iteration + fillRayPoolFromCamera<<>>(renderCam->resolution, (float)iterations, cam, traceDepth, cudaimage, cudageoms, + numberOfGeoms, cudarays); + + int numberOfRays = (int)renderCam->resolution.x * (int)renderCam->resolution.y; + + //std::cout<<"StreamCompaction: "; + + int linearTileSize = tileSize*tileSize; + for(int i=0; i < MAX_RECURSION_DEPTH && numberOfRays > 0; i++) { - raytraceRay<<>>(renderCam->resolution, (float)iterations, (float)bounce, cam, traceDepth, cudaimage, cudageoms, numberOfGeoms, cudamaterials, numberOfMaterials); +#if COMPACTION + dim3 linearGridSize((int)ceil(numberOfRays*1.0f/linearTileSize),1,1); + raytraceRay<<>>(renderCam->resolution, (float)iterations, cam, traceDepth+i, + cudaimage, cudageoms, numberOfGeoms, cudamaterials, numberOfMaterials, + i%2==0?cudarays : cudarays2, + numberOfRays); + checkCUDAError("Ray Trace Failed!"); + + numberOfRays = streamCompactRayPool( i%2==0? cudarays : cudarays2, + i%2==0? cudarays2 : cudarays, + numberOfRays); +#else + //kernel launches + raytraceRay<<>>(renderCam->resolution, (float)iterations, cam, traceDepth+i, cudaimage, cudageoms, numberOfGeoms, cudamaterials, numberOfMaterials, cudarays); +#endif } + + //getchar(); + +#if ACCUMULATION + combineIntoAccumulatorImage<<>>(renderCam->resolution, (float)iterations, cudaimage, accumulatorImage); + sendImageToPBO<<>>(PBOpos, renderCam->resolution, accumulatorImage, (float)iterations); +#else sendImageToPBO<<>>(PBOpos, renderCam->resolution, cudaimage); +#endif - //retrieve image from GPU - cudaMemcpy( renderCam->image, cudaimage, (int)renderCam->resolution.x*(int)renderCam->resolution.y*sizeof(glm::vec3), cudaMemcpyDeviceToHost); + //retrieve image from GPU for sending to bmp file + if(singleFrameMode) +#if ACCUMULATION + cudaMemcpy( renderCam->image, accumulatorImage, (int)renderCam->resolution.x*(int)renderCam->resolution.y*sizeof(glm::vec3), cudaMemcpyDeviceToHost); +#else + cudaMemcpy( renderCam->image, cudaimage, (int)renderCam->resolution.x*(int)renderCam->resolution.y*sizeof(glm::vec3), cudaMemcpyDeviceToHost); +#endif //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 + cudaFree( cudarays ); + delete geomList; +#if COMPACTION + cudaFree( cudarays2); +#endif + // make certain the kernel has completed cudaThreadSynchronize(); checkCUDAError("Kernel failed!"); } + +//Clear AccumulatorImage. For an interactive application, this needs to be called everytime the camera moves or the scene changes +void cudaClearAccumulatorImage(camera *renderCam) +{ + // 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))); + + clearImage<<>>(renderCam->resolution, accumulatorImage); +} + +//Allocate Memory For AccumulatorImage +void cudaAllocateAccumulatorImage(camera *renderCam) +{ + cudaMalloc((void**)&accumulatorImage, (int)renderCam->resolution.x*(int)renderCam->resolution.y*sizeof(glm::vec3)); + cudaClearAccumulatorImage(renderCam); +} + +//Free memory of the accumulator image +void cudaFreeAccumulatorImage() +{ + cudaFree(accumulatorImage); +} \ No newline at end of file diff --git a/src/raytraceKernel.h b/src/raytraceKernel.h index 5fcf5a3..3673a66 100755 --- a/src/raytraceKernel.h +++ b/src/raytraceKernel.h @@ -20,6 +20,13 @@ #include #endif -void cudaRaytraceCore(uchar4* pos, camera* renderCam, int frame, int iterations, material* materials, int numberOfMaterials, geom* geoms, int numberOfGeoms); +extern glm::vec3* accumulatorImage; + +void cudaAllocateAccumulatorImage(camera *renderCam); +void cudaFreeAccumulatorImage(); +void cudaClearAccumulatorImage(camera *renderCam); +void cudaRaytraceCore(uchar4* pos, camera* renderCam, int frame, int iterations, material* materials, int numberOfMaterials, geom* geoms, int numberOfGeoms, cameraData liveCamera); + +__host__ __device__ glm::vec3 generateRandomNumberFromThread(glm::vec2 resolution, float time, int x, int y); #endif diff --git a/src/scene.cpp b/src/scene.cpp index 415d627..93fcb97 100755 --- a/src/scene.cpp +++ b/src/scene.cpp @@ -162,6 +162,8 @@ int scene::loadCamera(){ vector positions; vector views; vector ups; + vector apertures; + vector focusPlanes; while (!line.empty() && fp_in.good()){ //check frame number @@ -172,7 +174,7 @@ int scene::loadCamera(){ } //load camera properties - for(int i=0; i<3; i++){ + for(int i=0; i<5; i++){ //glm::vec3 translation; glm::vec3 rotation; glm::vec3 scale; utilityCore::safeGetline(fp_in,line); tokens = utilityCore::tokenizeString(line); @@ -183,6 +185,12 @@ int scene::loadCamera(){ }else if(strcmp(tokens[0].c_str(), "UP")==0){ ups.push_back(glm::vec3(atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str()))); } + else if(strcmp(tokens[0].c_str(), "APERTURE")==0){ + apertures.push_back(atof(tokens[1].c_str())); + }else if(strcmp(tokens[0].c_str(), "FOCUSPLANE")==0){ + focusPlanes.push_back(atof(tokens[1].c_str())); + } + } frameCount++; @@ -194,10 +202,14 @@ int scene::loadCamera(){ newCamera.positions = new glm::vec3[frameCount]; newCamera.views = new glm::vec3[frameCount]; newCamera.ups = new glm::vec3[frameCount]; + newCamera.apertures = new float[frameCount]; + newCamera.focusPlanes = new float[frameCount]; for(int i=0; i