diff --git a/PROJ1_WIN/565Raytracer.suo b/PROJ1_WIN/565Raytracer.suo new file mode 100644 index 0000000..a578286 Binary files /dev/null and b/PROJ1_WIN/565Raytracer.suo differ diff --git a/PROJ1_WIN/565Raytracer/Release/565Raytracer.Build.CppClean.log b/PROJ1_WIN/565Raytracer/Release/565Raytracer.Build.CppClean.log new file mode 100644 index 0000000..ae09fe5 --- /dev/null +++ b/PROJ1_WIN/565Raytracer/Release/565Raytracer.Build.CppClean.log @@ -0,0 +1,21 @@ +H:\CIS565 GPU\GITHUB\PROJECT2-PATHTRACER\PROJ1_WIN\565RAYTRACER\RELEASE\565RAYTRACER.EXE.INTERMEDIATE.MANIFEST +H:\CIS565 GPU\GitHub\Project2-Pathtracer\PROJ1_WIN\565Raytracer\Release\cl.command.1.tlog +H:\CIS565 GPU\GitHub\Project2-Pathtracer\PROJ1_WIN\565Raytracer\Release\CL.read.1.tlog +H:\CIS565 GPU\GitHub\Project2-Pathtracer\PROJ1_WIN\565Raytracer\Release\CL.write.1.tlog +H:\CIS565 GPU\GITHUB\PROJECT2-PATHTRACER\PROJ1_WIN\565RAYTRACER\RELEASE\GLSLUTILITY.OBJ +H:\CIS565 GPU\GITHUB\PROJECT2-PATHTRACER\PROJ1_WIN\565RAYTRACER\RELEASE\IMAGE.OBJ +H:\CIS565 GPU\GitHub\Project2-Pathtracer\PROJ1_WIN\565Raytracer\Release\link.command.1.tlog +H:\CIS565 GPU\GitHub\Project2-Pathtracer\PROJ1_WIN\565Raytracer\Release\link.read.1.tlog +H:\CIS565 GPU\GitHub\Project2-Pathtracer\PROJ1_WIN\565Raytracer\Release\link.write.1.tlog +H:\CIS565 GPU\GITHUB\PROJECT2-PATHTRACER\PROJ1_WIN\565RAYTRACER\RELEASE\MAIN.OBJ +H:\CIS565 GPU\GitHub\Project2-Pathtracer\PROJ1_WIN\565Raytracer\Release\mt.command.1.tlog +H:\CIS565 GPU\GitHub\Project2-Pathtracer\PROJ1_WIN\565Raytracer\Release\mt.read.1.tlog +H:\CIS565 GPU\GitHub\Project2-Pathtracer\PROJ1_WIN\565Raytracer\Release\mt.write.1.tlog +H:\CIS565 GPU\GitHub\Project2-Pathtracer\PROJ1_WIN\565Raytracer\Release\raytraceKernel.cu.cache +H:\CIS565 GPU\GITHUB\PROJECT2-PATHTRACER\PROJ1_WIN\565RAYTRACER\RELEASE\SCENE.OBJ +H:\CIS565 GPU\GITHUB\PROJECT2-PATHTRACER\PROJ1_WIN\565RAYTRACER\RELEASE\STB_IMAGE.OBJ +H:\CIS565 GPU\GITHUB\PROJECT2-PATHTRACER\PROJ1_WIN\565RAYTRACER\RELEASE\STB_IMAGE_WRITE.OBJ +H:\CIS565 GPU\GITHUB\PROJECT2-PATHTRACER\PROJ1_WIN\565RAYTRACER\RELEASE\UTILITIES.OBJ +H:\CIS565 GPU\GITHUB\PROJECT2-PATHTRACER\PROJ1_WIN\565RAYTRACER\RELEASE\VC100.PDB +H:\CIS565 GPU\GITHUB\PROJECT2-PATHTRACER\PROJ1_WIN\RELEASE\565RAYTRACER.EXE +H:\CIS565 GPU\GITHUB\PROJECT2-PATHTRACER\PROJ1_WIN\RELEASE\565RAYTRACER.PDB diff --git a/PROJ1_WIN/565Raytracer/Release/565Raytracer.log b/PROJ1_WIN/565Raytracer/Release/565Raytracer.log new file mode 100644 index 0000000..b8cc86b --- /dev/null +++ b/PROJ1_WIN/565Raytracer/Release/565Raytracer.log @@ -0,0 +1,15 @@ +Build started 10/12/2012 11:27:58 PM. + 1>Project "H:\CIS565 GPU\GitHub\Project2-Pathtracer\PROJ1_WIN\565Raytracer\565Raytracer.vcxproj" on node 2 (clean target(s)). + 1>_PrepareForClean: + Deleting file "Release\565Raytracer.lastbuildstate". + CudaClean: + cmd.exe /C "C:\Users\ntndeyx\AppData\Local\Temp\f4793ac191f940d7af98e53079fd1efa.cmd" + "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.0\bin\nvcc.exe" -ccbin "S:\VS2010\VC\bin" -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" -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.0\include" --keep-dir "Release" -maxrregcount=0 --machine 32 --compile -Xcompiler "/EHsc /nologo /Od /Zi /MD " -o "H:\CIS565 GPU\GitHub\Project2-Pathtracer\PROJ1_WIN\565Raytracer\Win32/Release/raytraceKernel.cu.obj" "H:\CIS565 GPU\GitHub\Project2-Pathtracer\src\raytraceKernel.cu" -clean + + H:\CIS565 GPU\GitHub\Project2-Pathtracer\PROJ1_WIN\565Raytracer>"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.0\bin\nvcc.exe" -ccbin "S:\VS2010\VC\bin" -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" -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.0\include" --keep-dir "Release" -maxrregcount=0 --machine 32 --compile -Xcompiler "/EHsc /nologo /Od /Zi /MD " -o "H:\CIS565 GPU\GitHub\Project2-Pathtracer\PROJ1_WIN\565Raytracer\Win32/Release/raytraceKernel.cu.obj" "H:\CIS565 GPU\GitHub\Project2-Pathtracer\src\raytraceKernel.cu" -clean + Deleting file "Release\raytraceKernel.cu.deps". + 1>Done Building Project "H:\CIS565 GPU\GitHub\Project2-Pathtracer\PROJ1_WIN\565Raytracer\565Raytracer.vcxproj" (clean target(s)). + +Build succeeded. + +Time Elapsed 00:00:02.59 diff --git a/PROJ1_WIN/565Raytracer/vc100.pdb b/PROJ1_WIN/565Raytracer/vc100.pdb new file mode 100644 index 0000000..bab4482 Binary files /dev/null and b/PROJ1_WIN/565Raytracer/vc100.pdb differ diff --git a/README.md b/README.md index cbc1dce..9c658cb 100755 --- a/README.md +++ b/README.md @@ -5,33 +5,21 @@ Fall 2012 ------------------------------------------------------------------------------- Due Friday, 10/12/2012 ------------------------------------------------------------------------------- - -------------------------------------------------------------------------------- -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 Karl as soon as possible. - ------------------------------------------------------------------------------- -INTRODUCTION: +Implemented Features: ------------------------------------------------------------------------------- -In this project, you will extend your raytracer from Project 1 into a full CUDA based global illumination pathtracer. - -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. +* 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 string compaction (see the Physically-based shading and pathtracing lecture slides from 09/24 if you don't know what this refers to) +* Perfect specular reflection +* Depth of field +* Refraction( still working on it) ------------------------------------------------------------------------------- -CONTENTS: +BLOG ------------------------------------------------------------------------------- -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. - -The Windows and OSX versions of the project build and run exactly the same way as in Project0 and Project1. +http://cghuyue.blogspot.com/ ------------------------------------------------------------------------------- REQUIREMENTS: @@ -71,45 +59,7 @@ Alternatively, implementing just one of the following features can satisfy the " 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! -------------------------------------------------------------------------------- -NOTES ON GLM: -------------------------------------------------------------------------------- -This project uses GLM, the GL Math library, for linear algebra. You need to know two important points on how GLM is used in this project: - -* In this project, indices in GLM vectors (such as vec3, vec4), are accessed via swizzling. So, instead of v[0], v.x is used, and instead of v[1], v.y is used, and so on and so forth. -* GLM Matrix operations work fine on NVIDIA Fermi cards and later, but pre-Fermi cards do not play nice with GLM matrices. As such, in this project, GLM matrices are replaced with a custom matrix struct, called a cudaMat4, found in cudaMat4.h. A custom function for multiplying glm::vec4s and cudaMat4s is provided as multiplyMV() in intersections.h. - -------------------------------------------------------------------------------- -BLOG -------------------------------------------------------------------------------- -As mentioned in class, all students should have student blogs detailing progress on projects. If you already have a blog, you can use it; otherwise, please create a blog using www.blogger.com or any other tool, such as www.wordpress.org. Blog posts on your project are due on the SAME DAY as the project, and should include: -* A brief description of the project and the specific features you implemented. -* A link to your github repo if the code is open source. -* At least one screenshot of your project running. -* A 30 second or longer video of your project running. To create the video use http://www.microsoft.com/expression/products/Encoder4_Overview.aspx -------------------------------------------------------------------------------- -THIRD PARTY CODE POLICY -------------------------------------------------------------------------------- -* Use of any third-party code must be approved by asking on Piazza. If it is approved, all students are welcome to use it. Generally, we approve use of third-party code that is not a core part of the project. For example, for the ray tracer, we would approve using a third-party library for loading models, but would not approve copying and pasting a CUDA function for doing refraction. -* Third-party code must be credited in README.md. -* Using third-party code without its approval, including using another student's code, is an academic integrity violation, and will result in you receiving an F for the semester. -------------------------------------------------------------------------------- -SELF-GRADING -------------------------------------------------------------------------------- -* On the submission date, email your grade, on a scale of 0 to 100, to Karl, yiningli@seas.upenn.edu, with a one paragraph explanation. Be concise and realistic. Recall that we reserve 30 points as a sanity check to adjust your grade. Your actual grade will be (0.7 * your grade) + (0.3 * our grade). We hope to only use this in extreme cases when your grade does not realistically reflect your work - it is either too high or too low. In most cases, we plan to give you the exact grade you suggest. -* Projects are not weighted evenly, e.g., Project 0 doesn't count as much as the path tracer. We will determine the weighting at the end of the semester based on the size of each project. - -------------------------------------------------------------------------------- -SUBMISSION -------------------------------------------------------------------------------- -As with the previous project, you should fork this project and work inside of your fork. Upon completion, commit your finished project back to your fork, and make a pull request to the master repository. -You should include a README.md file in the root directory detailing the following -* A brief description of the project and specific features you implemented -* At least one screenshot of your project running, and at least one screenshot of the final rendered output of your pathtracer -* Instructions for building and running your project if they differ from the base code -* A link to your blog post detailing the project -* A list of all third-party code used \ No newline at end of file diff --git a/renders/huyue_reflection (1).jpg b/renders/huyue_reflection (1).jpg new file mode 100644 index 0000000..417791a Binary files /dev/null and b/renders/huyue_reflection (1).jpg differ diff --git a/renders/huyue_screenshot.jpg b/renders/huyue_screenshot.jpg new file mode 100644 index 0000000..380ffd1 Binary files /dev/null and b/renders/huyue_screenshot.jpg differ diff --git a/scenes/sampleScene.txt b/scenes/sampleScene.txt index 936135b..5866115 100755 --- a/scenes/sampleScene.txt +++ b/scenes/sampleScene.txt @@ -38,7 +38,7 @@ MATERIAL 3 //red glossy RGB .63 .06 .04 SPECEX 0 SPECRGB 1 1 1 -REFL 0 +REFL 1 REFR 0 REFRIOR 2 SCATTER 0 @@ -50,8 +50,8 @@ MATERIAL 4 //white glossy RGB 1 1 1 SPECEX 0 SPECRGB 1 1 1 -REFL 0 -REFR 0 +REFL 1 +REFR 1 REFRIOR 2 SCATTER 0 ABSCOEFF 0 0 0 @@ -74,7 +74,7 @@ MATERIAL 6 //green glossy RGB .15 .48 .09 SPECEX 0 SPECRGB 1 1 1 -REFL 0 +REFL 1 REFR 0 REFRIOR 2.6 SCATTER 0 diff --git a/src/glm/CMakeLists (2).txt b/src/glm/CMakeLists (2).txt new file mode 100644 index 0000000..d5ba209 --- /dev/null +++ b/src/glm/CMakeLists (2).txt @@ -0,0 +1,43 @@ +set(NAME glm) + +file(GLOB ROOT_SOURCE *.cpp) +file(GLOB ROOT_INLINE *.inl) +file(GLOB ROOT_HEADER *.hpp) + +file(GLOB_RECURSE CORE_SOURCE ./core/*.cpp) +file(GLOB_RECURSE CORE_INLINE ./core/*.inl) +file(GLOB_RECURSE CORE_HEADER ./core/*.hpp) + +file(GLOB_RECURSE GTC_SOURCE ./gtc/*.cpp) +file(GLOB_RECURSE GTC_INLINE ./gtc/*.inl) +file(GLOB_RECURSE GTC_HEADER ./gtc/*.hpp) + +file(GLOB_RECURSE GTX_SOURCE ./gtx/*.cpp) +file(GLOB_RECURSE GTX_INLINE ./gtx/*.inl) +file(GLOB_RECURSE GTX_HEADER ./gtx/*.hpp) + +file(GLOB_RECURSE VIRTREV_SOURCE ./virtrev/*.cpp) +file(GLOB_RECURSE VIRTREV_INLINE ./virtrev/*.inl) +file(GLOB_RECURSE VIRTREV_HEADER ./virtrev/*.hpp) + +source_group("Core Files" FILES ${CORE_SOURCE}) +source_group("Core Files" FILES ${CORE_INLINE}) +source_group("Core Files" FILES ${CORE_HEADER}) +source_group("GTC Files" FILES ${GTC_SOURCE}) +source_group("GTC Files" FILES ${GTC_INLINE}) +source_group("GTC Files" FILES ${GTC_HEADER}) +source_group("GTX Files" FILES ${GTX_SOURCE}) +source_group("GTX Files" FILES ${GTX_INLINE}) +source_group("GTX Files" FILES ${GTX_HEADER}) +source_group("VIRTREV Files" FILES ${VIRTREV_SOURCE}) +source_group("VIRTREV Files" FILES ${VIRTREV_INLINE}) +source_group("VIRTREV Files" FILES ${VIRTREV_HEADER}) + +include_directories(${CMAKE_CURRENT_SOURCE_DIR}/..) + +add_executable(${NAME} + ${ROOT_SOURCE} ${ROOT_INLINE} ${ROOT_HEADER} + ${CORE_SOURCE} ${CORE_INLINE} ${CORE_HEADER} + ${GTC_SOURCE} ${GTC_INLINE} ${GTC_HEADER} + ${GTX_SOURCE} ${GTX_INLINE} ${GTX_HEADER} + ${VIRTREV_SOURCE} ${VIRTREV_INLINE} ${VIRTREV_HEADER}) diff --git a/src/interactions.h b/src/interactions.h index e18cfff..5da7422 100755 --- a/src/interactions.h +++ b/src/interactions.h @@ -40,22 +40,44 @@ __host__ __device__ bool calculateScatterAndAbsorption(ray& r, float& depth, Ab //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); + float n = incidentIOR / transmittedIOR; + float cosAngle = glm::dot(normal, -incident); + float k = 1.0f - ( n*n * (1.0 - cosAngle * cosAngle)); + glm::vec3 transmission = n * incident + (n * cosAngle - sqrt(k)) * normal; + + return glm::normalize(transmission); + //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); +__host__ __device__ glm::vec3 calculateReflectionDirection(glm::vec3 normal, glm::vec3 incident) +{ + //nothing fancy here, R=2(N.L)N+L + //return glm::vec3(0,0,0); + glm::vec3 rlect = incident - 2.0f * glm::dot(incident,normal) * normal ; + return rlect; } //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; + //fresnel.reflectionCoefficient = 1; + //fresnel.transmissionCoefficient = 0; + + float cosAngleIn = glm::dot(incident, normal); + float n1 = incidentIOR, n2 = transmittedIOR; + float tmpA = (n1*n1) / (n2*n2) * (1.0f - cosAngleIn * cosAngleIn); + float cosAngleTran = sqrt(1.0f - tmpA); + + float Rs = ( (n1*cosAngleIn - n2*cosAngleTran) / ( n1*cosAngleIn + n2*cosAngleTran)) * + ( (n1*cosAngleIn - n2*cosAngleTran) / ( n1*cosAngleIn + n2*cosAngleTran)) ; + float Rp = ( (n1*cosAngleTran - n2*cosAngleIn) / ( n2*cosAngleIn + n1*cosAngleTran)) * + ( (n1*cosAngleTran - n2*cosAngleIn) / ( n2*cosAngleIn + n1*cosAngleTran)) ; + + fresnel.reflectionCoefficient = (Rs + Rp) / 2.0f; + fresnel.transmissionCoefficient = 1 - fresnel.reflectionCoefficient; + + return fresnel; } //LOOK: This function demonstrates cosine weighted random direction generation in a sphere! @@ -89,17 +111,43 @@ __host__ __device__ glm::vec3 calculateRandomDirectionInHemisphere(glm::vec3 nor //TODO: IMPLEMENT THIS FUNCTION //Now that you know how cosine weighted direction generation works, try implementing non-cosine (uniform) weighted random direction generation. //This should be much easier than if you had to implement calculateRandomDirectionInHemisphere. -__host__ __device__ glm::vec3 getRandomDirectionInSphere(float xi1, float xi2) { - return glm::vec3(0,0,0); +__host__ __device__ glm::vec3 getRandomDirectionInSphere(float xi1, float xi2) +{//??????????what is xi1 and xi2??????????? + //assume they are random numbers in [0,1]; + float theta = TWO_PI * xi1; + float phi = acos( 2*xi2 - 1); + + return ( glm::vec3( cos(theta) * sin(phi), sin(theta) * sin(phi), cos(phi) ) ); + } //TODO (PARTIALLY OPTIONAL): IMPLEMENT THIS FUNCTION //returns 0 if diffuse scatter, 1 if reflected, 2 if transmitted. __host__ __device__ int calculateBSDF(ray& r, glm::vec3 intersect, glm::vec3 normal, glm::vec3 emittedColor, AbsorptionAndScatteringProperties& currentAbsorptionAndScattering, - glm::vec3& color, glm::vec3& unabsorbedColor, material m){ - - return 1; + glm::vec3& color, glm::vec3& unabsorbedColor, material m, float xi1, float xi2,int closestGeomIndex){ + //Reflective material + if(m.hasReflective > 0) + { + r.continueFlag = true; + r.direction = calculateReflectionDirection(normal, r.direction); + r.origin = intersect + r.direction * 0.01f; + color = m.color; + } + + //Refractive material--still working on it + else if(m.hasRefractive > 0) ; + + //Diffuse material + else + { + r.continueFlag = true; + r.direction = calculateRandomDirectionInHemisphere(glm::normalize(normal), xi1, xi2); + r.origin = intersect + r.direction * 0.01f; + color = m.color; + } + + return 0; }; #endif diff --git a/src/intersections.h b/src/intersections.h index 714e918..908218c 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,104 @@ __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 botbcklft = glm::vec3(-0.5, -0.5, -0.5);//bottom back left + glm::vec3 topfrtrt = glm::vec3(0.5, 0.5, 0.5);//top front right + + 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; + //Ray-Box Intersection: http://www.siggraph.org/education/materials/HyperGraph/raytrace/rtinter3.htm + float t = 0; + float Tnear = -10000.0;//-infinity + float Tfar = 10000.0;//infinity + + if(rd.x == 0) + {//parallel to the X planes + if(ro.x < botbcklft.x || ro.x > topfrtrt.x) return -1; + } + else + { + float T1 = (botbcklft.x - ro.x) / rd.x; + float T2 = (topfrtrt.x - ro.x) / rd.x; + float tmp; + if(T1 > T2) + {//swap (T1, T2) /* since T1 intersection with near plane */ + tmp =T1; + T1 = T2; + T2 = tmp; + } + if(T1 > Tnear) Tnear = T1; + if(T2 < Tfar) Tfar = T2; + if(Tnear > Tfar) return -1; + if(Tfar < 0) return -1; + } + + if(rd.y == 0) + {//parallel to the Y planes + if(ro.y < botbcklft.y || ro.y > topfrtrt.y) return -1; + } + else + { + float T1 = (botbcklft.y - ro.y) / rd.y; + float T2 = (topfrtrt.y - ro.y) / rd.y; + float tmp; + if(T1 > T2) + {//swap (T1, T2) /* since T1 intersection with near plane */ + tmp =T1; + T1 = T2; + T2 = tmp; + } + if(T1 > Tnear) Tnear = T1; + if(T2 < Tfar) Tfar = T2; + if(Tnear > Tfar) return -1; + if(Tfar < 0) return -1; + } + + if(rd.z == 0) + {//parallel to the Z planes + if(ro.z < botbcklft.z || ro.z > topfrtrt.z) return -1; + } + else + { + float T1 = (botbcklft.z - ro.z) / rd.z; + float T2 = (topfrtrt.z - ro.z) / rd.z; + float tmp; + if(T1 > T2) + {//swap (T1, T2) /* since T1 intersection with near plane */ + tmp =T1; + T1 = T2; + T2 = tmp; + } + if(T1 > Tnear) Tnear = T1; + if(T2 < Tfar) Tfar = T2; + if(Tnear > Tfar) return -1; + if(Tfar < 0) return -1; + } + + if( Tnear >= -10000.0f ) t = Tnear;// t is the intersection distance + else return -1; + + glm::vec3 realIntersectionPoint = multiplyMV(box.transform, glm::vec4(getPointOnRay(rt, t), 1.0)); + glm::vec3 realOrigin = multiplyMV(box.transform, glm::vec4(0,0,0,1)); + intersectionPoint = realIntersectionPoint; + //normal = glm::normalize(realIntersectionPoint - realOrigin); + + glm::vec3 inP = rt.origin + rt.direction * Tnear; + if( fabs(inP.x - 0.5) < 0.0005 ) normal = glm::vec3(1.0f,.0f,.0f); + else if( fabs(inP.x + 0.5) < 0.0005 ) normal = glm::vec3(-1.0f,.0f,.0f); + if( fabs(inP.y - 0.5) < 0.0005 ) normal = glm::vec3(.0f,1.0f,.0f); + else if( fabs(inP.y + 0.5) < 0.0005 ) normal = glm::vec3(.0f,-1.0f,.0f); + if( fabs(inP.z - 0.5) < 0.0005 ) normal = glm::vec3(.0f,.0f,1.0f); + else if( fabs(inP.z + 0.5) < 0.0005 ) normal = glm::vec3(.0f,.0f,-1.0f); + + + normal = multiplyMV( box.transform, glm::vec4(normal, .0f) ); + normal = glm::normalize(normal); + + return glm::length(r.origin - realIntersectionPoint); } //LOOK: Here's an intersection test example from a sphere. Now you just need to figure out cube and, optionally, triangle. @@ -261,28 +266,20 @@ __host__ __device__ glm::vec3 getRandomPointOnCube(staticGeom cube, float random } +//TODO: IMPLEMENT THIS FUNCTION //Generates a random point on a given sphere -__host__ __device__ glm::vec3 getRandomPointOnSphere(staticGeom sphere, float randomSeed){ - - float radius=.5f; - thrust::default_random_engine rng(hash(randomSeed)); - thrust::uniform_real_distribution u01(0,1); - thrust::uniform_real_distribution u02(-0.5,0.5); - - glm::vec3 point = glm::vec3(0,0,0); - float x=(float)u02(rng); - float y=(float)u02(rng); - float z=0; - float russianRoulette = (float)u01(rng); - if(russianRoulette<0.5){ - z=(float)sqrt(radius*radius-x*x-y*y); - }else - z=-(float)sqrt(radius*radius-x*x-y*y); - - point=glm::vec3(x,y,z); - glm::vec3 randPoint = multiplyMV(sphere.transform, glm::vec4(point,1.0f)); - - return randPoint; +__host__ __device__ glm::vec3 getRandomPointOnSphere(staticGeom sphere, float randomSeed) +{ + thrust::default_random_engine rng(hash(randomSeed)); + thrust::uniform_real_distribution u(0,1); + thrust::uniform_real_distribution v(0,1); + + float theta = TWO_PI * u(rng); + float phi = acos( 2*v(rng) - 1 ); + glm::vec3 point = glm::vec3( cos(theta) * sin(phi), sin(theta) * sin(phi), cos(phi) ) +sphere.translation; + glm::vec3 randPoint = multiplyMV( sphere.transform, glm::vec4(point,1.0f) ); + + return randPoint; } #endif \ No newline at end of file diff --git a/src/main.cpp b/src/main.cpp index 4e94892..ea5362c 100755 --- a/src/main.cpp +++ b/src/main.cpp @@ -77,6 +77,8 @@ int main(int argc, char** argv){ glUseProgram(passthroughProgram); glActiveTexture(GL_TEXTURE0); + + #ifdef __APPLE__ // send into GLFW main loop while(1){ @@ -224,6 +226,24 @@ void runCuda(){ switch (key) { case(27): + case 'w': + renderCam->positions->z += 0.5f; + break; + case 's': + renderCam->positions->z -= 0.5f; + break; + case 'a': + renderCam->positions->x += 0.5f; + break; + case 'd': + renderCam->positions->x -= 0.5f; + break; + case 'u': + renderCam->positions->y += 0.5f; + break; + case 'i': + renderCam->positions->y -= 0.5f; + break; exit(1); break; } diff --git a/src/raytraceKernel.cu b/src/raytraceKernel.cu index d473c89..fd704c3 100755 --- a/src/raytraceKernel.cu +++ b/src/raytraceKernel.cu @@ -4,18 +4,24 @@ // Rob Farber for CUDA-GL interop, from CUDA Supercomputing For The Masses: http://www.drdobbs.com/architecture-and-design/cuda-supercomputing-for-the-masses-part/222600097 // Peter Kutz and Yining Karl Li's GPU Pathtracer: http://gpupathtracer.blogspot.com/ // Yining Karl Li's TAKUA Render, a massively parallel pathtracing renderer: http://www.yiningkarlli.com +#include +#include +#include +#include #include #include #include #include "sceneStructs.h" #include +#include "glm/glm.hpp" #include "utilities.h" #include "raytraceKernel.h" #include "intersections.h" #include "interactions.h" #include -#include "glm/glm.hpp" + +#define traceDepth 10 void checkCUDAError(const char *msg) { cudaError_t err = cudaGetLastError(); @@ -36,42 +42,28 @@ __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 and caches the result. "First bounce cache, second bounce thrash!" -__host__ __device__ ray raycastFromCameraKernel(glm::vec2 resolution, float time, int x, int y, glm::vec3 eye, glm::vec3 view, glm::vec3 up, glm::vec2 fov){ - - int index = x + (y * resolution.x); - - thrust::default_random_engine rng(hash(index*time)); - thrust::uniform_real_distribution u01(0,1); - - //standard camera raycast stuff - glm::vec3 E = eye; - glm::vec3 C = view; - glm::vec3 U = up; - float fovx = fov.x; - float fovy = fov.y; - - float CD = glm::length(C); - - glm::vec3 A = glm::cross(C, U); - glm::vec3 B = glm::cross(A, C); - glm::vec3 M = E+C; - glm::vec3 H = (A*float(CD*tan(fovx*(PI/180))))/float(glm::length(A)); - glm::vec3 V = (B*float(CD*tan(-fovy*(PI/180))))/float(glm::length(B)); - - float sx = (x)/(resolution.x-1); - float sy = (y)/(resolution.y-1); - - glm::vec3 P = M + (((2*sx)-1)*H) + (((2*sy)-1)*V); - glm::vec3 PmE = P-E; - glm::vec3 R = E + (float(200)*(PmE))/float(glm::length(PmE)); - - glm::vec3 direction = glm::normalize(R); - //major performance cliff at this point, TODO: find out why! - ray r; - r.origin = eye; - r.direction = direction; - return r; +//TODO: IMPLEMENT THIS FUNCTION +//Function that does the initial raycast from the camera +__host__ __device__ ray raycastFromCameraKernel(glm::vec2 resolution, float time, int x, int y, glm::vec3 eye, glm::vec3 view, glm::vec3 up, glm::vec2 fov) +{ + ray r; + r.origin = eye; + + glm::vec3 AVEC,BVEC,MVEC,HVEC,VVEC,Ppoint;//from CIS560 + float Sx = x / (resolution.x ); + float Sy = y / (resolution.y ); + + AVEC = glm::cross(view, up);//view is the CVEC, up is UVEC + BVEC = glm::cross(AVEC, view); + MVEC = eye + view;//Midpoint of screen + HVEC = view.length() * tan(fov.x) * glm::normalize(AVEC); + VVEC = view.length() * tan(fov.y) * glm::normalize(BVEC); + Ppoint = MVEC + ( 2*Sx - 1 ) * HVEC + ( 2*Sy -1 ) * VVEC; + + r.direction = glm::normalize(Ppoint - eye); + r.continueFlag = true; + + return r; } //Kernel that blacks out a given image buffer @@ -118,52 +110,153 @@ __global__ void sendImageToPBO(uchar4* PBOpos, glm::vec2 resolution, glm::vec3* } } +//generate rays for further ray tracing +__global__ void generateRay(ray *rays, cameraData cam, float iter, float focus) +{ + int x = (blockIdx.x * blockDim.x) + threadIdx.x; + int y = (blockIdx.y * blockDim.y) + threadIdx.y; + int index = x + (y * cam.resolution.x); + + //for anti-aliasing + thrust::default_random_engine rng( hash(index * iter) ); + thrust::uniform_real_distribution X(-0.5, 0.5); + float u = X(rng); + float v = X(rng); + + if(x <= cam.resolution.x && y <= cam.resolution.y) + { + //rays[index] = raycastFromCameraKernel(cam.resolution, 0.0f, x, y, cam.position, cam.view, cam.up, cam.fov); + // /////////////////////////////////////////////////////////////////////////////////////// + //Anti-aliasing + rays[index] = raycastFromCameraKernel(cam.resolution, 0.0f, x+u, y+v, cam.position, cam.view, cam.up, cam.fov); + rays[index].pixelId = index; + } + + ////////////////////////////////////////////////////////////////////////////////////////////// + //DEPTH FIELD + glm::vec3 FOC = rays[index].origin + rays[index].direction * focus; + thrust::uniform_real_distribution Y(-0.4, 0.4); + float offsetX = Y(rng), offsetY = Y(rng); + rays[index].origin += glm::vec3(offsetX, offsetY, 0.0f); + rays[index].direction = glm::normalize(FOC - rays[index].origin); + ///////////////////////////////////////////////////////////////////////////////////////// + +} + +//This function returns the closest Geometry's ID +__host__ __device__ int getClosestGeom(staticGeom* geoms, int numberOfGeoms, ray &r, glm::vec3 &intersecP, glm::vec3 &normal) +{ + int closestGeomIndex = -1; + float min_d = 10000; + float d;//distance to the nearest geometry + + for( int i = 0; i < numberOfGeoms; ++i ) + {//this loop find out the closest object and light source + if(geoms[i].type == SPHERE) d = sphereIntersectionTest(geoms[i],r,intersecP,normal); + else if(geoms[i].type == CUBE ) d = boxIntersectionTest(geoms[i],r,intersecP,normal); + + if( d > 0 && d < min_d ) + { + min_d = d; + closestGeomIndex = i; + } + } + return closestGeomIndex; +} +//////////////////////////////////////////////////////////////////////////////////////////////////////////////////// +__global__ void Accumulate(glm::vec2 resolution, glm::vec3* colors, float iterations, glm::vec3* current) +//__global__ void Accumulate(glm::vec2 resolution, glm::vec3* colors, float iterations) +{//accumulate color + int x = (blockIdx.x * blockDim.x) + threadIdx.x; + int y = (blockIdx.y * blockDim.y) + threadIdx.y; + int index = x + (y * resolution.x); + colors[index] = ( colors[index] + current[index] * ( iterations- 1 ) ) / iterations; // / iterations; +} + //TODO: IMPLEMENT THIS FUNCTION //Core raytracer kernel __global__ void raytraceRay(glm::vec2 resolution, float time, cameraData cam, int rayDepth, glm::vec3* colors, - staticGeom* geoms, int numberOfGeoms, material* materials, int numberOfMaterials){ + staticGeom* geoms, int numberOfGeoms, material* materials, ray* rays, int numOfRays)//Added cudaMaterial +{ + int x = (blockIdx.x * blockDim.x) + threadIdx.x; + int y = (blockIdx.y * blockDim.y) + threadIdx.y; + int index = x + (y * resolution.x);//pixel - int x = (blockIdx.x * blockDim.x) + threadIdx.x; - int y = (blockIdx.y * blockDim.y) + threadIdx.y; - int index = x + (y * resolution.x); + if(x <= cam.resolution.x && y <= cam.resolution.y) + {//ray r = raycastFromCameraKernel(resolution,time,x,y,cam.position,cam.view,cam.up,cam.fov); + ray r = rays[index]; - ray r = raycastFromCameraKernel(resolution, time, x, y, cam.position, cam.view, cam.up, cam.fov); - - if((x<=resolution.x && y<=resolution.y)){ - - float MAX_DEPTH = 100000000000000000; - float depth = MAX_DEPTH; - - for(int i=0; i-EPSILON){ - MAX_DEPTH = depth; - colors[index] = materials[geoms[i].materialid].color; - } - } + glm::vec3 colResult; + int hitCounter = 0; + if(rayDepth == 1 ) colors[r.pixelId] = glm:: vec3(1,1,1); + if( index > numOfRays || r.continueFlag==false ) return; + glm::vec3 intersecP, norm; + float d = -1;//distance to intersection + float min_d = 1000000.0;//the distance to closest object + glm::vec3 lightPos; + float lightEmi; + glm::vec3 lightCol; + int lightGeoIndex=-1;//store the index of the light source, Only one currently + float amibient = 0.8, diffuse = 0.8, specular = 0.1; + + glm::vec3 emittedColor; + glm::vec3 unabsorbedColor; + AbsorptionAndScatteringProperties AbsorpASP; - //colors[index] = generateRandomNumberFromThread(resolution, time, x, y); - } -} + int closestGeomIndex = getClosestGeom(geoms, numberOfGeoms, r, intersecP, norm); + + //intersection occurred + if(closestGeomIndex >= 0 ) + { + material geoMat = materials[ geoms[closestGeomIndex].materialid ]; + glm::vec3 geoCol; + //colors[r.pixelId] = glm::vec3(1,1,1); + + // Light + if( geoMat.emittance > 0) + { + colResult = geoMat.color * geoMat.emittance;// the light color + r.continueFlag = false;//don't need to keep going + } + + //Non - Light Objects + else + { + geoCol = geoMat.color; + thrust::default_random_engine rng(hash(index*time*rayDepth)); + thrust::uniform_real_distribution X1(0, 1); + float xi1 = X1(rng); + float xi2 = X1(rng); + calculateBSDF(r,intersecP, norm, emittedColor, AbsorpASP, colResult, unabsorbedColor, geoMat, xi1, xi2,closestGeomIndex); + }//END of the intersected object is not light source + + }//END of intersect with object + //NO INTERSECTION + else + { + colResult = glm::vec3(0,0,0); + r.continueFlag = false; + } + + colors[r.pixelId] *= colResult; + rays[index] = r; + + if(rayDepth + 1 > traceDepth && r.continueFlag) + { + colors[r.pixelId] = glm:: vec3(0, 0, 0); + r.continueFlag = false; + } + }//this is for the if(x <= cam.resolution.x && y <= cam.resolution.y) +} //TODO: FINISH THIS FUNCTION // Wrapper for the __global__ call that sets up the kernel calls and does a ton of memory management void cudaRaytraceCore(uchar4* PBOpos, camera* renderCam, int frame, int iterations, material* materials, int numberOfMaterials, geom* geoms, int numberOfGeoms){ - int traceDepth = 1; //determines how many bounces the raytracer traces + //int traceDepth = 1; //determines how many bounces the raytracer traces // set up crucial magic int tileSize = 8; @@ -172,9 +265,14 @@ void cudaRaytraceCore(uchar4* PBOpos, camera* renderCam, int frame, int iteratio //send image to GPU glm::vec3* cudaimage = NULL; + glm::vec3 *current = NULL; + cudaMalloc((void**)&cudaimage, (int)renderCam->resolution.x*(int)renderCam->resolution.y*sizeof(glm::vec3)); cudaMemcpy( cudaimage, renderCam->image, (int)renderCam->resolution.x*(int)renderCam->resolution.y*sizeof(glm::vec3), cudaMemcpyHostToDevice); - + + cudaMalloc((void**)¤t, (int)renderCam->resolution.x*(int)renderCam->resolution.y*sizeof(glm::vec3)); + cudaMemcpy( current, renderCam->image, (int)renderCam->resolution.x*(int)renderCam->resolution.y*sizeof(glm::vec3), cudaMemcpyHostToDevice); + //package geometry and materials and sent to GPU staticGeom* geomList = new staticGeom[numberOfGeoms]; for(int i=0; iups[frame]; cam.fov = renderCam->fov; - //kernel launches - raytraceRay<<>>(renderCam->resolution, (float)iterations, cam, traceDepth, cudaimage, cudageoms, numberOfGeoms, cudamaterials, - numberOfMaterials); + //Package rays + int numOfRays = cam.resolution.x * cam.resolution.y; + ray *rays = new ray[numOfRays]; + ray *cudarays = NULL; + cudaMalloc((void**)&cudarays, numOfRays * sizeof(ray)); + cudaMemcpy(cudarays, rays, numOfRays * sizeof(ray), cudaMemcpyHostToDevice); + //////////////////////////////////////////////Focused on the green sphere's Z + float focus =17; + generateRay<<>>(cudarays,cam,iterations,focus); + //kernel launches + //traceDepth = 10; + + dim3 blocksPerGrid = fullBlocksPerGrid; + + for(int i=1; i < traceDepth + 1; ++i) + { + //raytraceRay(resolution, time, cam, rayDepth, colors, geoms, numberOfGeoms, materials, rays, numOfRays) + raytraceRay<<>>(renderCam->resolution, (float)iterations, cam, i, cudaimage, cudageoms, numberOfGeoms,cudaMaterials,cudarays,numOfRays); + thrust::device_ptr firstRays_ptr(cudarays); + thrust::device_ptr lastRays_ptr = thrust::remove_if( firstRays_ptr, firstRays_ptr + numOfRays, rayContinueFalse()); + numOfRays = lastRays_ptr.get() - firstRays_ptr.get(); + blocksPerGrid = dim3( (int)ceil(renderCam->resolution.x/tileSize), (int)ceil(renderCam->resolution.x)/tileSize); + } + Accumulate<<>>( renderCam->resolution, cudaimage, iterations, current); sendImageToPBO<<>>(PBOpos, renderCam->resolution, cudaimage); //retrieve image from GPU cudaMemcpy( renderCam->image, cudaimage, (int)renderCam->resolution.x*(int)renderCam->resolution.y*sizeof(glm::vec3), cudaMemcpyDeviceToHost); //free up stuff, or else we'll leak memory like a madman - cudaFree( cudaimage ); - cudaFree( cudageoms ); - cudaFree( cudamaterials ); - delete [] geomList; + cudaFree( cudaimage ); + cudaFree( current ); + cudaFree( cudageoms ); + cudaFree( cudaMaterials ); + cudaFree( cudarays ); + + + delete geomList; + delete materialsList; + delete rays; // make certain the kernel has completed cudaThreadSynchronize(); diff --git a/src/sceneStructs.h b/src/sceneStructs.h index b10f1cf..c9a3b5c 100755 --- a/src/sceneStructs.h +++ b/src/sceneStructs.h @@ -16,6 +16,14 @@ enum GEOMTYPE{ SPHERE, CUBE, MESH }; struct ray { glm::vec3 origin; glm::vec3 direction; + bool continueFlag; + int pixelId; + bool insideFlag; + int curObjId; +}; + +struct rayContinueFalse{ + __host__ __device__ bool operator()(ray r) {return !r.continueFlag;} }; struct geom { diff --git a/src/utilities.h b/src/utilities.h index 5842c33..cb0e89d 100755 --- a/src/utilities.h +++ b/src/utilities.h @@ -17,13 +17,13 @@ #include #include "cudaMat4.h" -const float PI =3.1415926535897932384626422832795028841971; -const float TWO_PI =6.2831853071795864769252867665590057683943; -const float SQRT_OF_ONE_THIRD =0.5773502691896257645091487805019574556476; -const float E =2.7182818284590452353602874713526624977572; -const float EPSILON =.000000001; -const float ZERO_ABSORPTION_EPSILON =0.00001; -const float RAY_BIAS_AMOUNT =0.0002; +#define PI 3.1415926535897932384626422832795028841971 +#define TWO_PI 6.2831853071795864769252867665590057683943 +#define SQRT_OF_ONE_THIRD 0.5773502691896257645091487805019574556476 +#define E 2.7182818284590452353602874713526624977572 +#define EPSILON .000000001 +#define ZERO_ABSORPTION_EPSILON 0.00001 +#define RAY_BIAS_AMOUNT 0.0002 namespace utilityCore { extern float clamp(float f, float min, float max);