diff --git a/PROJ1_WIN/565Raytracer/565Raytracer.vcxproj b/PROJ1_WIN/565Raytracer/565Raytracer.vcxproj
index fcc853d..36c7c6d 100755
--- a/PROJ1_WIN/565Raytracer/565Raytracer.vcxproj
+++ b/PROJ1_WIN/565Raytracer/565Raytracer.vcxproj
@@ -34,7 +34,9 @@
-
+
+ false
+
{FF21CA49-522E-4E86-B508-EE515B248FC4}
@@ -56,7 +58,7 @@
-
+
@@ -78,7 +80,7 @@
Level3
Disabled
WIN32;_DEBUG;_CONSOLE;%(PreprocessorDefinitions)
- C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.0\include;C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.0\C\common\inc;../shared/glew/include;../shared/freeglut/include;%(AdditionalIncludeDirectories)
+ 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
@@ -91,7 +93,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
+ 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
@@ -103,7 +105,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
@@ -116,11 +118,11 @@
$(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
-
+
\ No newline at end of file
diff --git a/PROJ1_WIN/565Raytracer/565Raytracer.vcxproj.filters b/PROJ1_WIN/565Raytracer/565Raytracer.vcxproj.filters
index d49ad9c..0e47e14 100755
--- a/PROJ1_WIN/565Raytracer/565Raytracer.vcxproj.filters
+++ b/PROJ1_WIN/565Raytracer/565Raytracer.vcxproj.filters
@@ -1,42 +1,80 @@
-
-
-
-
-
stb_image
stb_image
+
+ Source
+
+
+ Source
+
+
+ Source
+
+
+ Source
+
+
+ Source
+
-
-
-
-
-
-
-
-
-
-
stb_image
stb_image
-
-
-
+
+ Header
+
+
+ Header
+
+
+ Header
+
+
+ Header
+
+
+ Header
+
+
+ Header
+
+
+ Header
+
+
+ Header
+
+
+ Header
+
+
+ Header
+
{011aa553-95e8-4e59-b7ff-1bb89aebe21d}
+
+ {b6af916f-7350-4c57-81e3-fd26535c112b}
+
+
+ {8b6462cf-94cd-461a-94b7-6d56a9a513e8}
+
+
+
+
+ Source
+
\ No newline at end of file
diff --git a/PROJ1_WIN/565Raytracer/565Raytracer.vcxproj.user b/PROJ1_WIN/565Raytracer/565Raytracer.vcxproj.user
index d7ca222..269a8da 100755
--- a/PROJ1_WIN/565Raytracer/565Raytracer.vcxproj.user
+++ b/PROJ1_WIN/565Raytracer/565Raytracer.vcxproj.user
@@ -1,7 +1,7 @@
- scene="../../scenes/sampleScene.txt"
+ scene="../../scenes/portalScene.txt"
WindowsLocalDebugger
\ No newline at end of file
diff --git a/PROJ1_WIN/Release/565Raytracer.exe b/PROJ1_WIN/Release/565Raytracer.exe
new file mode 100644
index 0000000..aeca672
Binary files /dev/null and b/PROJ1_WIN/Release/565Raytracer.exe differ
diff --git a/PROJ1_WIN/Release/565Raytracer.pdb b/PROJ1_WIN/Release/565Raytracer.pdb
new file mode 100644
index 0000000..34e3081
Binary files /dev/null and b/PROJ1_WIN/Release/565Raytracer.pdb differ
diff --git a/PROJ1_WIN/Release/freeglut.dll b/PROJ1_WIN/Release/freeglut.dll
new file mode 100644
index 0000000..4ec8893
Binary files /dev/null and b/PROJ1_WIN/Release/freeglut.dll differ
diff --git a/PROJ1_WIN/Release/glew32.dll b/PROJ1_WIN/Release/glew32.dll
new file mode 100644
index 0000000..fc4d8d8
Binary files /dev/null and b/PROJ1_WIN/Release/glew32.dll differ
diff --git a/PROJ1_WIN/Release/glut32.dll b/PROJ1_WIN/Release/glut32.dll
new file mode 100644
index 0000000..106646f
Binary files /dev/null and b/PROJ1_WIN/Release/glut32.dll differ
diff --git a/README.md b/README.md
index cbc1dce..deef41c 100755
--- a/README.md
+++ b/README.md
@@ -7,109 +7,17 @@ Due Friday, 10/12/2012
-------------------------------------------------------------------------------
-------------------------------------------------------------------------------
-NOTE:
+Path Tracer
-------------------------------------------------------------------------------
-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.
+For this project I've extended the ray tracer into a path tracer. All of the base requirements have been met and I have added in camera controls and a fancy new primitive: the portal!
--------------------------------------------------------------------------------
-INTRODUCTION:
--------------------------------------------------------------------------------
-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.
-
--------------------------------------------------------------------------------
-CONTENTS:
--------------------------------------------------------------------------------
-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.
-
--------------------------------------------------------------------------------
-REQUIREMENTS:
--------------------------------------------------------------------------------
-In this project, you are given code for:
-
-* 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!
-
-* 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
-
-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.
-
-* 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
+Camera controls demoed in the video.
-Alternatively, implementing just one of the following features can satisfy the "pick two" feature requirement, since these are correspondingly more difficult problems:
+Portals, like the ones in the scene are comprised of a planar colliding surface defined by a Julia set the parameters of which are defined in the scene file (circles are a special case of the Julia set with C set to 0). Each portal is linked to another portal, allowing for a ray to enter portal A and exit portal B. I have added the functionality that input and output portals need not be the same size, this allows for scaling effects that make objects appear larger or smaller or in some way warped. Portals may map to any other portal. This means that you can have one portal be the output of many different inputs, i.e. portals A, B, and C all lead to portal D. I use this characteristic to create the back lighting on the circular and fractal portals in the scene by adding a duplicate of the portals behind the main viewing ones that all map to a portal placed outside of the bounding box facing a light source.
-* Physically based subsurface scattering and transmission
-* Implement and integrate your own stackless KD-Tree from scratch.
-* Displacement mapping
-* Deformational motion blur
-
-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.
+As can be seen by comparing this video to the last, I have upgraded my hardware to an Nvidia GTX 680 and have seen substantial speed boosts.
-------------------------------------------------------------------------------
-SELF-GRADING
+BLOG:
-------------------------------------------------------------------------------
-* 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
+http://liamboone.blogspot.com/2012/10/project-2-pathtracer.html
\ No newline at end of file
diff --git a/scenes/portalScene.txt b/scenes/portalScene.txt
new file mode 100644
index 0000000..42ab647
--- /dev/null
+++ b/scenes/portalScene.txt
@@ -0,0 +1,241 @@
+MATERIAL 0 //white diffuse
+RGB 0.85 0.9 0.75
+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 0.6 0.7 0
+SPECEX 0
+SPECRGB 1 1 1
+REFL 0
+REFR 0
+REFRIOR 2
+SCATTER 0
+ABSCOEFF 0 0 0
+RSCTCOEFF 0
+EMITTANCE 0
+
+MATERIAL 5 //glass
+RGB 0 0 0
+SPECEX 0
+SPECRGB 1 1 1
+REFL 0
+REFR 1
+REFRIOR 2.2
+SCATTER 0
+ABSCOEFF .02 5.1 5.7
+RSCTCOEFF 13
+EMITTANCE 0
+
+MATERIAL 6 //green glossy
+RGB .15 .09 .48
+SPECEX 0
+SPECRGB 1 1 1
+REFL 0
+REFR 0
+REFRIOR 2.6
+SCATTER 0
+ABSCOEFF 0 0 0
+RSCTCOEFF 0
+EMITTANCE 0
+
+MATERIAL 7 //light
+RGB 1 1 1
+SPECEX 0
+SPECRGB 0 0 0
+REFL 0
+REFR 0
+REFRIOR 0
+SCATTER 0
+ABSCOEFF 0 0 0
+RSCTCOEFF 0
+EMITTANCE 1
+
+MATERIAL 8 //light
+RGB 1 1 1
+SPECEX 0
+SPECRGB 0 0 0
+REFL 0
+REFR 0
+REFRIOR 0
+SCATTER 0
+ABSCOEFF 0 0 0
+RSCTCOEFF 0
+EMITTANCE 15
+
+CAMERA
+RES 800 800
+FOVY 45
+ITERATIONS 5000
+FILE renders/sampleScene.bmp
+frame 0
+EYE 0 4.5 4
+VIEW 0 0 -1
+UP 0 1 0
+frame 1
+EYE 0 4.5 12
+VIEW 0 0 -1
+UP 0 1 0
+
+OBJECT 0
+cube
+material 0
+frame 0
+TRANS 0 0 0
+ROTAT 0 0 90
+SCALE .01 10 10
+
+OBJECT 1
+cube
+material 0
+frame 0
+TRANS 0 5 -5
+ROTAT 0 90 0
+SCALE .01 10 10
+
+OBJECT 2
+cube
+material 0
+frame 0
+TRANS 0 10 0
+ROTAT 0 0 90
+SCALE .01 10 10
+
+OBJECT 3
+cube
+material 1
+frame 0
+TRANS -5 5 0
+ROTAT 0 0 90
+SCALE 10 .01 10
+
+OBJECT 4
+cube
+material 2
+frame 0
+TRANS 5 5 0
+ROTAT 0 0 90
+SCALE 10 .01 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
+portal 0 0 //0.2524 0.5577
+material 9
+frame 0
+TRANS -4.8 5 0
+ROTAT 0 90 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
+portal 0 0 //0.2524 0.5577
+material 7
+frame 0
+TRANS 4.8 5 0
+ROTAT 0 90 0
+SCALE 2.5 2.5 2.5
+
+OBJECT 10
+portal 0 0 //0.2524 0.5577
+material 12
+frame 0
+TRANS -4.9 5 0
+ROTAT 0 90 0
+SCALE 3 3 3
+
+OBJECT 11
+portal 0 0 //0.2524 0.5577
+material 12
+frame 0
+TRANS 4.9 5 0
+ROTAT 0 270 0
+SCALE 2.5 2.5 2.5
+
+OBJECT 12
+portal 0 0 //0.2524 0.5577
+material 12
+frame 0
+TRANS 0 10.5 0
+ROTAT 90 0 0
+SCALE 0.5 0.5 0.5
+
+OBJECT 13
+cube
+material 6
+frame 0
+TRANS 0 5 5
+ROTAT 0 90 0
+SCALE .01 10 10
+
+OBJECT 14
+portal -1.1566 -0.2642
+material 12
+frame 0
+TRANS 0 5 -4.8
+ROTAT 0 0 0
+SCALE 3 3 3
\ No newline at end of file
diff --git a/scenes/sampleScene.txt b/scenes/sampleScene.txt
index 936135b..32518dc 100755
--- a/scenes/sampleScene.txt
+++ b/scenes/sampleScene.txt
@@ -161,8 +161,8 @@ cube
material 1
frame 0
TRANS -5 5 0
-ROTAT 0 0 0
-SCALE .01 10 10
+ROTAT 0 0 90
+SCALE 10 .01 10
frame 1
TRANS -5 5 0
ROTAT 0 0 0
@@ -173,8 +173,8 @@ cube
material 2
frame 0
TRANS 5 5 0
-ROTAT 0 0 0
-SCALE .01 10 10
+ROTAT 0 0 90
+SCALE 10 .01 10
frame 1
TRANS 5 5 0
ROTAT 0 0 0
@@ -226,4 +226,4 @@ SCALE .3 3 3
frame 1
TRANS 0 10 0
ROTAT 0 0 90
-SCALE .3 3 3
\ No newline at end of file
+SCALE .3 3 3
diff --git a/src/interactions.h b/src/interactions.h
index e18cfff..a764a99 100755
--- a/src/interactions.h
+++ b/src/interactions.h
@@ -63,7 +63,7 @@ __host__ __device__ glm::vec3 calculateRandomDirectionInHemisphere(glm::vec3 nor
//crucial difference between this and calculateRandomDirectionInSphere: THIS IS COSINE WEIGHTED!
- float up = sqrt(xi1); // cos(theta)
+ float up = xi1; // cos(theta)
float over = sqrt(1 - up * up); // sin(theta)
float around = xi2 * TWO_PI;
@@ -83,7 +83,6 @@ __host__ __device__ glm::vec3 calculateRandomDirectionInHemisphere(glm::vec3 nor
glm::vec3 perpendicularDirection2 = glm::normalize(glm::cross(normal, perpendicularDirection1));
return ( up * normal ) + ( cos(around) * over * perpendicularDirection1 ) + ( sin(around) * over * perpendicularDirection2 );
-
}
//TODO: IMPLEMENT THIS FUNCTION
diff --git a/src/intersections.h b/src/intersections.h
index 714e918..0d34ee7 100755
--- a/src/intersections.h
+++ b/src/intersections.h
@@ -69,9 +69,98 @@ __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));
}
+__host__ __device__ float portalIntersectionTest(staticGeom inP, staticGeom outP, ray r, glm::vec3& intersectionPoint, glm::vec3& normal, float x1 )
+{
+ glm::vec3 ro = multiplyMV(inP.inverseTransform, glm::vec4(r.origin,1.0f));
+ glm::vec3 rd = glm::normalize(multiplyMV(inP.inverseTransform, glm::vec4(r.direction,0.0f)));
+
+ ray rt; rt.origin = ro; rt.direction = rd;
+
+ float t = -rt.origin.z / rt.direction.z;
+
+ glm::vec3 p = rt.origin + t * rt.direction;
+
+ if( ( p.x*p.x + p.y*p.y ) > 4 || t < 0 )
+ {
+ return -1;
+ }
+
+ glm::vec2 juliaC = inP.julia;
+ glm::vec2 juliaZ = p.swizzle( glm::comp::X, glm::comp::Y );
+
+ int i = 0;
+
+ for( ; i < 50; i++ )
+ {
+ float real = juliaZ.x*juliaZ.x - juliaZ.y*juliaZ.y;
+ float complex = 2*juliaZ.x*juliaZ.y;
+ juliaZ = glm::vec2(real, complex) + juliaC;
+ if( glm::length( juliaZ ) > 2 ) break;
+ }
+
+ float x2 = max( (float)(i-5)/45.0f, 0.0f );
+
+ if( x2*x2 < x1 ) return -1;
+
+ glm::vec3 realIntersectionPoint = multiplyMV(outP.transform, glm::vec4(getPointOnRay(rt, t+0.0002), 1.0));
+
+ intersectionPoint = realIntersectionPoint;
+ normal = glm::normalize( multiplyMV(outP.transform, glm::vec4(rt.direction, 0.0)) );
+
+ return glm::length(r.origin - multiplyMV(inP.transform, glm::vec4(getPointOnRay(rt, t), 1.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);
+ 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;
+
+ glm::vec3 tlfb = ( glm::vec3(-0.5,-0.5,-0.5) - ro ) / rd;
+ glm::vec3 trbt = ( glm::vec3(0.5,0.5,0.5) - ro ) / rd;
+
+ glm::vec3 tmin = glm::min( tlfb, trbt );
+ glm::vec3 tmax = glm::max( tlfb, trbt );
+
+ if( tmax.x < tmin.y || tmax.x < tmin.z || tmax.y < tmin.x ||
+ tmax.y < tmin.z || tmax.z < tmin.x || tmax.z < tmin.y )
+ {
+ return -1;
+ }
+
+ float t = ( tmin.x > tmin.y ? tmin.x : tmin.y );
+ t = ( t > tmin.z ? t : tmin.z );
+
+ if( t < 0 ) return -1;
+
+ glm::vec3 realNormal = glm::vec3( 1, 1, 1 );
+
+ //glm::vec4 point = glm::vec4( getPointOnRay(rt, t), 1.0f );
+ glm::vec4 point = glm::vec4( rt.origin + rt.direction*t, 1.0f );
+
+ if( fabs( point.x - 0.5 ) < 2e-4 )
+ realNormal = glm::vec3( 1, 0, 0 );
+ else if( fabs( point.x + 0.5 ) < 2e-4 )
+ realNormal = glm::vec3( -1, 0, 0 );
+ else if( fabs( point.y - 0.5 ) < 2e-4 )
+ realNormal = glm::vec3( 0, 1, 0 );
+ else if( fabs( point.y + 0.5 ) < 2e-4 )
+ realNormal = glm::vec3( 0, -1, 0 );
+ else if( fabs( point.z - 0.5 ) < 2e-4 )
+ realNormal = glm::vec3( 0, 0, 1 );
+ else if( fabs( point.z + 0.5 ) < 2e-4 )
+ realNormal = glm::vec3( 0, 0, -1 );
+
+ point = glm::vec4( getPointOnRay(rt, t), 1.0f );
+ glm::vec3 realIntersectionPoint = multiplyMV(box.transform, point);
+ glm::vec3 realOrigin = multiplyMV(box.transform, glm::vec4(0,0,0,1));
+
+ normal = glm::normalize(multiplyMV(box.transform, glm::vec4(realNormal,0.0f)));
+ intersectionPoint = realIntersectionPoint;
+
+ return glm::length(r.origin - realIntersectionPoint);
+ //return boxIntersectionTest(glm::vec3(-.5,-.5,-.5), glm::vec3(.5,.5,.5), box, r, intersectionPoint, normal);
}
//Cube intersection test, return -1 if no intersection, otherwise, distance to intersection
@@ -155,11 +244,11 @@ __host__ __device__ float boxIntersectionTest(glm::vec3 boxMin, glm::vec3 boxMa
currentNormal = glm::vec3(0,0,-1);
}
- intersectionPoint = multiplyMV(box.transform, glm::vec4(osintersect, 1.0));
+ intersectionPoint = multiplyMV(box.transform, glm::vec4(getPointOnRay(r, tmin), 1.0));
- normal = multiplyMV(box.transform, glm::vec4(currentNormal,0.0));
+ normal = glm::normalize(multiplyMV(box.transform, glm::vec4(currentNormal,0.0)));
return glm::length(intersectionPoint-ro.origin);
}
diff --git a/src/main.cpp b/src/main.cpp
index 4e94892..d35c792 100755
--- a/src/main.cpp
+++ b/src/main.cpp
@@ -100,6 +100,8 @@ int main(int argc, char** argv){
//---------RUNTIME STUFF---------
//-------------------------------
+int traceDepth = 2;
+
void runCuda(){
// Map OpenGL buffer object for writing from CUDA on a single GPU
@@ -123,7 +125,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(), traceDepth );
// unmap buffer object
cudaGLUnmapBufferObject(pbo);
@@ -197,7 +199,7 @@ void runCuda(){
}
#else
-
+
void display(){
runCuda();
@@ -220,12 +222,78 @@ void runCuda(){
void keyboard(unsigned char key, int x, int y)
{
- std::cout << key << std::endl;
+ glm::vec3 cross = glm::cross( renderCam->views[targetFrame], renderCam->ups[targetFrame] );
switch (key)
{
- case(27):
- exit(1);
- break;
+ case('w'):
+ renderCam->positions[targetFrame] += renderCam->views[targetFrame];
+ iterations = 0;
+ break;
+
+ case('s'):
+ renderCam->positions[targetFrame] -= renderCam->views[targetFrame];
+ iterations = 0;
+ break;
+
+ case('q'):
+ renderCam->positions[targetFrame] += renderCam->ups[targetFrame];
+ iterations = 0;
+ break;
+
+ case('z'):
+ renderCam->positions[targetFrame] -= renderCam->ups[targetFrame];
+ iterations = 0;
+ break;
+
+ case('a'):
+ renderCam->positions[targetFrame] += cross;
+ iterations = 0;
+ break;
+
+ case('d'):
+ renderCam->positions[targetFrame] -= cross;
+ iterations = 0;
+ break;
+
+ case('i'):
+ renderCam->views[targetFrame] += renderCam->ups[targetFrame] * 0.1f;
+ renderCam->views[targetFrame] = glm::normalize( renderCam->views[targetFrame] );
+ iterations = 0;
+ break;
+
+ case('k'):
+ renderCam->views[targetFrame] -= renderCam->ups[targetFrame] * 0.1f;
+ renderCam->views[targetFrame] = glm::normalize( renderCam->views[targetFrame] );
+ iterations = 0;
+ break;
+
+ case('j'):
+ renderCam->views[targetFrame] += cross * 0.1f;
+ renderCam->views[targetFrame] = glm::normalize( renderCam->views[targetFrame] );
+ iterations = 0;
+ break;
+
+ case('l'):
+ renderCam->views[targetFrame] -= cross * 0.1f;
+ renderCam->views[targetFrame] = glm::normalize( renderCam->views[targetFrame] );
+ iterations = 0;
+ break;
+
+ case(']'):
+ traceDepth = ( traceDepth + 1 ) % 50;
+ iterations = 0;
+ cout << traceDepth << endl;
+ break;
+
+ case('['):
+ traceDepth = ( traceDepth + 49 ) % 50;
+ iterations = 0;
+ cout << traceDepth << endl;
+ break;
+
+ case(27):
+ exit(1);
+ break;
}
}
diff --git a/src/raytraceKernel.cu b/src/raytraceKernel.cu
index d473c89..b3e2195 100755
--- a/src/raytraceKernel.cu
+++ b/src/raytraceKernel.cu
@@ -16,6 +16,10 @@
#include "interactions.h"
#include
#include "glm/glm.hpp"
+#include "cuda_runtime.h"
+#include
+
+int * numThreads;
void checkCUDAError(const char *msg) {
cudaError_t err = cudaGetLastError();
@@ -38,40 +42,41 @@ __host__ __device__ glm::vec3 generateRandomNumberFromThread(glm::vec2 resolutio
//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;
+
+ 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+(float)u01(rng))/(resolution.x);
+ float sy = (y+(float)u01(rng))/(resolution.y);
+
+ glm::vec3 P = M + (((2*sx)-1)*H) + (((2*sy)-1)*V);
+ glm::vec3 PmE = P-E;
+ glm::vec3 R = E + (float(200)*(PmE))/float(glm::length(PmE));
+
+ glm::vec3 direction = glm::normalize(R);
+ //major performance cliff at this point, TODO: find out why!
+ ray r;
+ r.origin = eye;
+ r.direction = direction;
+
+ return r;
}
//Kernel that blacks out a given image buffer
@@ -85,7 +90,7 @@ __global__ void clearImage(glm::vec2 resolution, glm::vec3* image){
}
//Kernel that writes the image to the OpenGL PBO directly.
-__global__ void sendImageToPBO(uchar4* PBOpos, glm::vec2 resolution, glm::vec3* image){
+__global__ void sendImageToPBO(uchar4* PBOpos, glm::vec2 resolution, glm::vec3* image, float iterations){
int x = (blockIdx.x * blockDim.x) + threadIdx.x;
int y = (blockIdx.y * blockDim.y) + threadIdx.y;
@@ -94,9 +99,9 @@ __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;
+ color.x = image[index].x*255.0/iterations;
+ color.y = image[index].y*255.0/iterations;
+ color.z = image[index].z*255.0/iterations;
if(color.x>255){
color.x = 255;
@@ -118,110 +123,285 @@ __global__ void sendImageToPBO(uchar4* PBOpos, glm::vec2 resolution, glm::vec3*
}
}
+__device__ void rayCast( ray r, staticGeom* geoms, int numberOfGeoms, bounce& b, int& obj, float seed )
+{
+ float MAX_DEPTH = 100000000000000000;
+ float depth = MAX_DEPTH;
+
+ obj = -1;
+
+ thrust::default_random_engine rng(hash(seed));
+ thrust::uniform_real_distribution u01(0,1);
+
+ glm::vec3 color( 0 );
+ glm::vec3 intersectionPoint( 0 );
+ glm::vec3 intersectionNormal( 0 );
+
+ for(int i=0; i-EPSILON)
+ {
+ obj = i;
+ MAX_DEPTH = depth;
+ b.matid = geoms[i].materialid;
+ b.geoid = i;
+ b.normal = intersectionNormal;
+ b.position = intersectionPoint;
+ }
+ }
+
+ b.incomingVector = r.direction;
+}
+
//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, int numberOfMaterials, bounce* bounces, int* numThreads)
+{
+ int x = (blockIdx.x * blockDim.x) + threadIdx.x;
+ int y = (blockIdx.y * blockDim.y) + threadIdx.y;
+ int index = x + (y * resolution.x);
- int x = (blockIdx.x * blockDim.x) + threadIdx.x;
- int y = (blockIdx.y * blockDim.y) + threadIdx.y;
- int index = x + (y * resolution.x);
+ if( index == 0 ) *numThreads = resolution.x*resolution.y;
- 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;
- }
- }
+ if((x<=resolution.x && y<=resolution.y))
+ {
+ int object;
+ ray r = raycastFromCameraKernel(resolution, time, x, y, cam.position, cam.view, cam.up, cam.fov);
+ rayCast( r, geoms, numberOfGeoms, bounces[index], object, time*index );
+ bounces[index].maxBounce = rayDepth;
+ bounces[index].pixel = index;
+
+ if( object == -1 )
+ {
+ bounces[index].color = glm::vec3( 0 );
+ bounces[index].count = rayDepth;
+ return;
+ }
+
+ bounces[index].color = glm::vec3( 1.0f ); //materials[geoms[object].materialid].color;
+ if( geoms[object].type != PORTAL )
+ bounces[index].count = ( materials[geoms[object].materialid].emittance > 0 ? rayDepth : 0 );
+ }
+}
+//naive stream compression with double buffering
+__global__ void compress(float its, int rayDepth, glm::vec3* colors, bounce* bounces, int* numThreads)
+{
+ int index = (blockIdx.x * blockDim.x) + threadIdx.x;
+ int idx;
- //colors[index] = generateRandomNumberFromThread(resolution, time, x, y);
- }
+ if( index < *numThreads )
+ {
+ int sum = 0;
+ if( bounces[index].count >= rayDepth )
+ {
+ idx = bounces[index].pixel;
+ colors[idx] += bounces[idx].color;
+ }
+ }
}
+__global__ void bounceAround(float time, cameraData cam, int rayDepth, glm::vec3* colors, staticGeom* geoms, int numberOfGeoms, material* materials, int numberOfMaterials, bounce* bounces, int* numThreads)
+{
+ int index = (blockIdx.x * blockDim.x) + threadIdx.x;
+ int reflection;
+
+ if( index < *numThreads )
+ {
+ if( bounces[index].count >= rayDepth )
+ return;
+
+ int timestamp = (int)time * index * (bounces[index].count+1);
+
+ thrust::default_random_engine rng(hash(timestamp));
+ thrust::uniform_real_distribution u01(0,1);
+ ray r;
+ float x0 = (float)u01(rng);
+ float x1 = (float)u01(rng);
+ if(geoms[bounces[index].geoid].type==PORTAL)
+ {
+ r.direction = bounces[index].normal;
+ reflection = 1;
+ }
+ else if( (float)u01(rng) < 0.5 || materials[bounces[index].matid].indexOfRefraction < EPSILON )
+ {
+ r.direction = calculateRandomDirectionInHemisphere( bounces[index].normal, x0, x1 );
+ reflection = 0;
+ }
+ else
+ {
+ reflection = 1;
+ r.direction = bounces[index].incomingVector - bounces[index].normal * ( 2 * glm::dot( bounces[index].incomingVector, bounces[index].normal ) );
+ }
+ r.direction = glm::normalize( r.direction );
+ r.origin = bounces[index].position;
+
+ int object;
+ material mat = materials[bounces[index].matid];
+ staticGeom geo = geoms[bounces[index].geoid];
+
+ rayCast( r, geoms, numberOfGeoms, bounces[index], object, timestamp );
+
+ if( object == -1 )
+ {
+ bounces[index].color = glm::vec3( 0 );
+ bounces[index].count = rayDepth;
+ return;
+ }
+
+ if( geo.type != PORTAL )
+ bounces[index].color *= ( mat.emittance > 0 ? mat.emittance : 1.0f );
+
+ if( !reflection )
+ {
+ bounces[index].color *= mat.color;
+ }
+
+ if( geo.type != PORTAL )
+ bounces[index].count = ( mat.emittance > 0 ? rayDepth : bounces[index].count+1 );
+
+ if( bounces[index].count >= rayDepth && mat.emittance <= EPSILON )
+ bounces[index].color *= 0;
+ }
+}
+
+struct is_done
+{
+ __host__ __device__
+ bool operator()( const bounce x )
+ {
+ return x.count < x.maxBounce;
+ }
+};
//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, int traceDepth){
- int traceDepth = 1; //determines how many bounces the raytracer traces
+ //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)));
+ // 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)));
+ dim3 threadsPerBounceBlock(tileSize*tileSize);
+ dim3 fullBounceBlocksPerGrid((int)ceil(float(renderCam->resolution.x*renderCam->resolution.y)/float(tileSize*tileSize)));
- //send image to GPU
- glm::vec3* cudaimage = NULL;
- cudaMalloc((void**)&cudaimage, (int)renderCam->resolution.x*(int)renderCam->resolution.y*sizeof(glm::vec3));
- cudaMemcpy( cudaimage, renderCam->image, (int)renderCam->resolution.x*(int)renderCam->resolution.y*sizeof(glm::vec3), cudaMemcpyHostToDevice);
+ //send image to GPU
+ glm::vec3* cudaimage = NULL;
+ cudaMalloc((void**)&cudaimage, (int)renderCam->resolution.x*(int)renderCam->resolution.y*sizeof(glm::vec3));
+ cudaMemcpy( cudaimage, renderCam->image, (int)renderCam->resolution.x*(int)renderCam->resolution.y*sizeof(glm::vec3), cudaMemcpyHostToDevice);
- //package geometry and materials and sent to GPU
- staticGeom* geomList = new staticGeom[numberOfGeoms];
- for(int i=0; iresolution;
- cam.position = renderCam->positions[frame];
- cam.view = renderCam->views[frame];
- cam.up = renderCam->ups[frame];
- cam.fov = renderCam->fov;
-
- //kernel launches
- raytraceRay<<>>(renderCam->resolution, (float)iterations, cam, traceDepth, cudaimage, cudageoms, numberOfGeoms, cudamaterials,
- numberOfMaterials);
+ material* cudamaterials = NULL;
+ cudaMalloc((void**)&cudamaterials, numberOfMaterials*sizeof(material));
+ cudaMemcpy( cudamaterials, materials, numberOfMaterials*sizeof(material), cudaMemcpyHostToDevice);
+
+ bounce* cudabounces = NULL;
+ bounce* Abounces = new bounce[ (int) (renderCam->resolution.x * renderCam->resolution.y) ];
+ bounce* Bbounces = new bounce[ (int) (renderCam->resolution.x * renderCam->resolution.y) ];
+ cudaMalloc( ( void ** ) &cudabounces, renderCam->resolution.x * renderCam->resolution.y * sizeof( bounce ) );
+
+ static bounce* firstbounces = NULL;
+ if( firstbounces == NULL ) cudaMalloc( ( void ** ) &firstbounces, renderCam->resolution.x * renderCam->resolution.y * sizeof( bounce ) );
+
+ checkCUDAError("1111111111111!");
+ //package camera
+ cameraData cam;
+ cam.resolution = renderCam->resolution;
+ cam.position = renderCam->positions[frame];
+ cam.view = renderCam->views[frame];
+ cam.up = renderCam->ups[frame];
+ cam.fov = renderCam->fov;
- sendImageToPBO<<>>(PBOpos, renderCam->resolution, cudaimage);
+ //kernel launches
+ if( iterations == 1 )
+ {
+ clearImage<<>>(renderCam->resolution, cudaimage);
+ cudaMalloc((void**)&(numThreads), sizeof(int));
+ }
+ raytraceRay<<>>(renderCam->resolution, (float)iterations, cam, traceDepth, cudaimage, cudageoms,
+ numberOfGeoms, cudamaterials, numberOfMaterials, firstbounces, numThreads);
+
+ checkCUDAError("2222222222222!");
+ cudaMemcpy( cudabounces, firstbounces, renderCam->resolution.x*renderCam->resolution.y*sizeof(bounce), cudaMemcpyDeviceToDevice);
+
+ checkCUDAError("3333333333333!");
+
+ int numT = 1;
- //retrieve image from GPU
- cudaMemcpy( renderCam->image, cudaimage, (int)renderCam->resolution.x*(int)renderCam->resolution.y*sizeof(glm::vec3), cudaMemcpyDeviceToHost);
+ for( int i = 0; i < traceDepth && numT > 0; i++ )
+ {
+ bounceAround<<>>((float)iterations, cam, traceDepth, cudaimage, cudageoms, numberOfGeoms, cudamaterials, numberOfMaterials, cudabounces, numThreads);
+ //compress<<>>((float)iterations, traceDepth, cudaimage, cudabounces, numThreads);
+ //cudaMemcpy( &numT, numThreads, sizeof( int ), cudaMemcpyDeviceToHost );
+ //cudaMemcpy( Abounces, cudabounces, numT * sizeof( bounce ), cudaMemcpyDeviceToHost );
+ //std::cout << numT << " " << Bbounces << " " << std::endl;
+ //bounce* qbounce = thrust::copy_if( Abounces, Abounces + numT, Bbounces, is_done() );
+ //numT = qbounce - Bbounces;
+ //std::cout << numT << " " << Bbounces << " " << qbounce << std::endl;
+ //system("pause");
+ //cudaMemcpy( numThreads, &numT, sizeof( int ), cudaMemcpyHostToDevice );
+ //cudaMemcpy( cudabounces, Bbounces, numT * sizeof( bounce ), cudaMemcpyHostToDevice );
+ }
- //free up stuff, or else we'll leak memory like a madman
- cudaFree( cudaimage );
- cudaFree( cudageoms );
- cudaFree( cudamaterials );
- delete [] geomList;
+ delete[] Abounces;
+ delete[] Bbounces;
+
+ checkCUDAError("In the Loop!");
+ compress<<>>((float)iterations, traceDepth, cudaimage, cudabounces, numThreads);
+
+ sendImageToPBO<<>>(PBOpos, renderCam->resolution, cudaimage, (float)iterations);
+
+ checkCUDAError("4444444444444!");
+ //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( cudabounces );
+ cudaFree( cudamaterials );
+ delete [] geomList;
- // make certain the kernel has completed
- cudaThreadSynchronize();
+ // make certain the kernel has completed
+ cudaThreadSynchronize();
- checkCUDAError("Kernel failed!");
+ checkCUDAError("Kernel failed!");
+ //system("pause");
}
diff --git a/src/raytraceKernel.h b/src/raytraceKernel.h
index 331e5ce..1518d76 100755
--- a/src/raytraceKernel.h
+++ b/src/raytraceKernel.h
@@ -15,6 +15,6 @@
#include "sceneStructs.h"
#include
-void cudaRaytraceCore(uchar4* pos, camera* renderCam, int frame, int iterations, material* materials, int numberOfMaterials, geom* geoms, int numberOfGeoms);
+void cudaRaytraceCore(uchar4* pos, camera* renderCam, int frame, int iterations, material* materials, int numberOfMaterials, geom* geoms, int numberOfGeoms, int traceDepth);
#endif
diff --git a/src/scene.cpp b/src/scene.cpp
index f0384b2..108d5aa 100755
--- a/src/scene.cpp
+++ b/src/scene.cpp
@@ -45,12 +45,18 @@ int scene::loadObject(string objectid){
//load object type
getline(fp_in,line);
if (!line.empty() && fp_in.good()){
+ vector tokens = utilityCore::tokenizeString(line);
if(strcmp(line.c_str(), "sphere")==0){
cout << "Creating new sphere..." << endl;
newObject.type = SPHERE;
}else if(strcmp(line.c_str(), "cube")==0){
cout << "Creating new cube..." << endl;
newObject.type = CUBE;
+ }else if(strcmp(tokens[0].c_str(), "portal")==0)
+ {
+ cout << "Creating new portal... Julia number: " << tokens[1].c_str() << " + j" << tokens[2].c_str() << endl;
+ newObject.type = PORTAL;
+ newObject.julia = glm::vec2( atof(tokens[1].c_str()), atof(tokens[2].c_str()) );
}else{
string objline = line;
string name;
diff --git a/src/sceneStructs.h b/src/sceneStructs.h
index b10f1cf..81ac2d6 100755
--- a/src/sceneStructs.h
+++ b/src/sceneStructs.h
@@ -11,7 +11,7 @@
#include
#include
-enum GEOMTYPE{ SPHERE, CUBE, MESH };
+enum GEOMTYPE{ SPHERE, CUBE, MESH, PORTAL };
struct ray {
glm::vec3 origin;
@@ -22,6 +22,7 @@ struct geom {
enum GEOMTYPE type;
int materialid;
int frames;
+ glm::vec2 julia;
glm::vec3* translations;
glm::vec3* rotations;
glm::vec3* scales;
@@ -32,6 +33,7 @@ struct geom {
struct staticGeom {
enum GEOMTYPE type;
int materialid;
+ glm::vec2 julia;
glm::vec3 translation;
glm::vec3 rotation;
glm::vec3 scale;
@@ -73,4 +75,15 @@ struct material{
float emittance;
};
+struct bounce{
+ glm::vec3 position;
+ glm::vec3 incomingVector;
+ glm::vec3 normal;
+ glm::vec3 color;
+ int maxBounce;
+ int pixel;
+ int count;
+ int matid;
+ int geoid;
+};
#endif //CUDASTRUCTS_H
diff --git a/src/utilities.h b/src/utilities.h
index 5842c33..922ff1e 100755
--- a/src/utilities.h
+++ b/src/utilities.h
@@ -17,13 +17,21 @@
#include
#include "cudaMat4.h"
+__device__ const float PI =3.1415926535897932384626422832795028841971;
+__device__ const float TWO_PI =6.2831853071795864769252867665590057683943;
+__device__ const float SQRT_OF_ONE_THIRD =0.5773502691896257645091487805019574556476;
+__device__ const float E =2.7182818284590452353602874713526624977572;
+__device__ const float EPSILON =.000000001;
+__device__ const float ZERO_ABSORPTION_EPSILON =0.00001;
+__device__ const float RAY_BIAS_AMOUNT =0.0002;
+/*
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;
+const float RAY_BIAS_AMOUNT =0.0002;*/
namespace utilityCore {
extern float clamp(float f, float min, float max);