diff --git a/README.md b/README.md index 605d944..4c7f1d8 100644 --- a/README.md +++ b/README.md @@ -3,274 +3,31 @@ CUDA Path Tracer **University of Pennsylvania, CIS 565: GPU Programming and Architecture, Project 3** -* (TODO) YOUR NAME HERE -* Tested on: (TODO) Windows 22, i7-2222 @ 2.22GHz 22GB, GTX 222 222MB (Moore 2222 Lab) +* Megan Moore +* Tested on: Windows 7, i7-4770 @ 3.40GHz 16GB (Moore 100 Lab C) -### (TODO: Your README) +![](img/motion_blur_1000.png "Motion blur after 1000 iterations") +* This is the final image of my GPU Path Tracing project. It shows a combination of different surfaces, motion blur, and antialiasing. The main objective of this project was to write a path tracer on the GPU that is much faster than a path tracer on the CPU. The reason we are able to do this is because a path tracer can be parallelized. Instead of using a recursive function in a large for loop that goes through each bounce of each ray, we can use a thread for each ray that will follow the rays through each bounce at the same time. Being able to calculate the color value of each pixel simultaneously, rather than one at a time, gives us a huge speed up in the run time of our path tracer. -*DO NOT* leave the README to the last minute! It is a crucial part of the -project, and we will not be able to grade you without a good README. +* Antialiasing: Antialiasing was done by jittering each ray as it was shot to a pixel. A random variable was created, and the ray was offset by that amount. This occured in every iteration on the GPU. As this occurred while each ray was being generated on the GPU, it did not cause any time to be added to the overall run time of the path tracer. The same would be true if it had been done on a CPU, as it is such a small amount of work. However, if I were to have done pixel sampling, which would have added the amount of rays that are shot for each pixel, this would have greatly increased the time, whether on a CPU or GPU. It would have caused less of a slow down on the GPU, however. It would have increased the amount of threads needed, but would still be faster than creating an even larger for loop on the CPU. -Instructions (delete me) -======================== +![](img/DOF1.png "Depth of Field") +* Depth of field: This was implemented by jittering the location of the aperatus, as the rays were being shot through it. Given a specific distance, we can choose which objects that are in focus and which are not. This was helpful to do on the GPU, because each ray could easily be jittered on the GPU. Again, like antialiasing, this did not cause any slow down on the overall run time. Doing this on the CPU would also not cause any slow down. It is a small calculation to change the original direction of the ray being shot at a pixel, so it does not require much memory or time. -This is due Thursday, September 24 evening at midnight. +![](img/specular1.png "Specular sphere") +* Specular: As you can see in the reference image at the top, and this image above, both perfectly specular and nonperfect specular surfaces have been implemented. The perfectly specular surface is a mirror and every ray that hits it will be reflected and bounce off of it, going off the surface at exactly the same angle it hit the surface. In a nonperfect specular surface, some rays will bounce off like it does on a mirror, and others will bounce off like it hit a diffuse surface, where each ray goes in a random direction. This allows there to be some reflection on the surface, but not quite as strong as a mirror, as seen in this image. If this were to be done on a CPU it would just allow the recursive function to continue following either the diffuse bounce or reflective bounce. However, no major speed ups or slow downs would occur based on implementing this surface on a CPU or GPU. -**Summary:** -In this project, you'll implement a CUDA-based path tracer capable of rendering -globally-illuminated images very quickly. -Since in this class we are concerned with working in GPU programming, -performance, and the generation of actual beautiful images (and not with -mundane programming tasks like I/O), this project includes base code for -loading a scene description file, described below, and various other things -that generally make up a framework for previewing and saving images. +![](img/refraction_correct3(glass).png "Refraction with an IOR of glass") +* Refraction: When a surface is refractive, it allows some rays to bounce off of it, while others go through the surface and come out at a different point. In order to implement refraction, we needed to use Schlick's approximation. This was used to calculate the probabilty that a ray would either bounce off or go through the surface. If a ray bounced off, it was treated like a mirror. If it went through the surface, it was refracted, and the exact point was then calculated and the ray would bounce out on the other side in a new direction. This allows for glass and water-like surfaces. The image above shows a surface with the index of refraction equal to that of glass. The image below shows one with an IOR equal to water. Again, there is no additional speed up or slow down from this feature. It occurs on the GPU and doesn't add any large memory needs. If this were to be done on a CPU, it would cause the recursive function to continue following the refractive bounce or reflective bounce. However, no major speed ups or slow downs would occur based on implementing this surface on a CPU or GPU. -The core renderer is left for you to implement. Finally, note that, while this -base code is meant to serve as a strong starting point for a CUDA path tracer, -you are not required to use it if you don't want to. You may also change any -part of the base code as you please. **This is YOUR project.** +![](img/refraction_correct2(water).png "Refraction with an IOR of water") -**Recommendation:** Every image you save should automatically get a different -filename. Don't delete all of them! For the benefit of your README, keep a -bunch of them around so you can pick a few to document your progress at the -end. +![](img/motion_blur1.png "Motion blur after 150 iterations") +* Motion Blur: Motion blur was implemented on the CPU rather than the GPU. This is because it needed to update the position of the moving object in every iteration. This meant the array the held all the objects needed to be updated. This could not be done on the GPU, individually for each ray. It needed to occur before the object array was passed onto the GPU. The way the motion blur was calculated, was that the object had a new property that was a vec3, that held the information of how far the object moved within the number of max iterations. Within each iteration, the object moved a portion of the to its new destination. As more iterations occur, the object moves to it's new position, with the rays hitting it in a different spot every iteration. Thus, creating a image that looks as if it has been in motion. This implementation may have caused a slight slow down, as the object array had to be updated every iteration, causing a memory transfer to occur every iteration, rather than just once at the beginning of the path trace. If this were to be done on the CPU, this memory transfor would not have to occur and it probably would not slow down the implementation all that much. -### Contents -* `src/` C++/CUDA source files. -* `scenes/` Example scene description files. -* `img/` Renders of example scene description files. - (These probably won't match precisely with yours.) -* `external/` Includes and static libraries for 3rd party libraries. - - -### Running the code - -The main function requires a scene description file. Call the program with -one as an argument: `cis565_path_tracer scene/sphere.txt`. -(In Visual Studio, `../scene/sphere.txt`.) - -If you are using Visual Studio, you can set this in the Debugging > Command -Arguments section in the Project properties. Make sure you get the path right - -read the console for errors. - -#### Controls - -* Esc to save an image and exit. -* Space to save an image. Watch the console for the output filename. -* W/A/S/D and R/F move the camera. Arrow keys rotate. - -## Requirements - -**Ask on the mailing list for clarifications.** - -In this project, you are given code for: - -* Loading and reading the scene description format -* Sphere and box intersection functions -* Support for saving images -* Working CUDA-GL interop for previewing your render while it's running -* A function which generates random screen noise (instead of an actual render). - -You will need to implement the following features: - -* Raycasting from the camera into the scene through an imaginary grid of pixels - (the screen) - * Implement antialiasing (by jittering rays within each pixel) -* Diffuse surfaces -* Perfectly specular-reflective (mirrored) surfaces - * See notes on diffuse/specular in `scatterRay` and on specular below -* Stream compaction optimization. You may use any of: - * Your global-memory work-efficient stream compaction implementation. - * A shared-memory work-efficient stream compaction (see below). - * `thrust::remove_if` or any of the other Thrust stream compaction functions. - -You are also required to implement at least 2 of the following features. -Please ask if you need good references (they will be added to this README -later on). If you find good references, share them! **Extra credit**: implement -more features on top of the 2 required ones, with point value up to +20/100 at -the grader's discretion (based on difficulty and coolness). - -* Work-efficient stream compaction using shared memory across multiple blocks - (See *GPU Gems 3* Chapter 39). -* These 2 smaller features: - * Refraction (e.g. glass/water) with Frensel effects using Schlick's - approximation or more accurate methods - * Physically-based depth-of-field (by jittering rays within an aperture) - * Recommended but not required: non-perfect specular surfaces -* Texture mapping -* Bump mapping -* Direct lighting (by taking a final ray directly to a random point on an - emissive object acting as a light source) -* Some method of defining object motion, and motion blur -* Subsurface scattering -* Arbitrary mesh loading and rendering (e.g. `obj` files). You can find these - online or export them from your favorite 3D modeling application. - With approval, you may use a third-party OBJ loading code to bring the data - into C++. - * You can use the triangle intersection function `glm::intersectRayTriangle`. - -This 'extra features' list is not comprehensive. If you have a particular idea -you would like to implement (e.g. acceleration structures, etc.), please -contact us first. - -For each extra feature, you must provide the following analysis: - -* Overview write-up of the feature -* Performance impact of the feature -* If you did something to accelerate the feature, what did you do and why? -* Compare your GPU version of the feature to a HYPOTHETICAL CPU version - (you don't have to implement it!) Does it benefit or suffer from being - implemented on the GPU? -* How might this feature be optimized beyond your current implementation? - -## Base Code Tour - -You'll be working in the following files. Look for important parts of the code: -search for `CHECKITOUT`. You'll have to implement parts labeled with `TODO`. -(But don't let these constrain you - you have free rein!) - -* `src/pathtrace.cu`: path tracing kernels, device functions, and calling code - * `pathtraceInit` initializes the path tracer state - it should copy - scene data (e.g. geometry, materials) from `Scene`. - * `pathtraceFree` frees memory allocated by `pathtraceInit` - * `pathtrace` performs one iteration of the rendering - it handles kernel - launches, memory copies, transferring some data, etc. - * See comments for a low-level path tracing recap. -* `src/intersections.h`: ray intersection functions - * `boxIntersectionTest` and `sphereIntersectionTest`, which take in a ray and - a geometry object and return various properties of the intersection. -* `src/interactions.h`: ray scattering functions - * `calculateRandomDirectionInHemisphere`: a cosine-weighted random direction - in a hemisphere. Needed for implementing diffuse surfaces. - * `scatterRay`: this function should perform all ray scattering, and will - call `calculateRandomDirectionInHemisphere`. See comments for details. -* `src/main.cpp`: you don't need to do anything here, but you can change the - program to save `.hdr` image files, if you want (for postprocessing). - -### Generating random numbers - -``` -thrust::default_random_engine rng(hash(index)); -thrust::uniform_real_distribution u01(0, 1); -float result = u01(rng); -``` - -There is a convenience function for generating a random engine using a -combination of index, iteration, and depth as the seed: - -``` -thrust::default_random_engine rng = random_engine(iter, index, depth); -``` - -### Specular Lighting - -In path tracing, like diffuse materials, specular materials are -simulated using a probability distribution instead computing the -strength of a ray bounce based on angles. - -Equations 7, 8, and 9 of -[GPU Gems 3 chapter 20](http://http.developer.nvidia.com/GPUGems3/gpugems3_ch20.html) -give the formulas for generating a random specular ray. (Note that -there is a typographical error: χ in the text = ξ in the formulas.) - -Also see the notes in `scatterRay` for probability splits between -diffuse/specular/other material types. - - -### Notes on GLM - -This project uses GLM for linear algebra. - -On NVIDIA cards pre-Fermi (pre-DX12), you may have issues with mat4-vec4 -multiplication. If you have one of these cards, be careful! If you have issues, -you might need to grab `cudamat4` and `multiplyMV` from the -[Fall 2014 project](https://github.com/CIS565-Fall-2014/Project3-Pathtracer). -Let us know if you need to do this. - -### Scene File Format - -This project uses a custom scene description format. Scene files are flat text -files that describe all geometry, materials, lights, cameras, and render -settings inside of the scene. Items in the format are delimited by new lines, -and comments can be added using C-style `// comments`. - -Materials are defined in the following fashion: - -* MATERIAL (material ID) //material header -* RGB (float r) (float g) (float b) //diffuse color -* SPECX (float specx) //specular exponent -* SPECRGB (float r) (float g) (float b) //specular color -* REFL (bool refl) //reflectivity flag, 0 for no, 1 for yes -* REFR (bool refr) //refractivity flag, 0 for no, 1 for yes -* REFRIOR (float ior) //index of refraction for Fresnel effects -* EMITTANCE (float emittance) //the emittance of the material. Anything >0 - makes the material a light source. - -Cameras are defined in the following fashion: - -* CAMERA //camera header -* RES (float x) (float y) //resolution -* FOVY (float fovy) //vertical field of view half-angle. the horizonal angle is calculated from this and the reslution -* ITERATIONS (float interations) //how many iterations to refine the image, - only relevant for supersampled antialiasing, depth of field, area lights, and - other distributed raytracing applications -* DEPTH (int depth) //maximum depth (number of times the path will bounce) -* FILE (string filename) //file to output render to upon completion -* EYE (float x) (float y) (float z) //camera's position in worldspace -* VIEW (float x) (float y) (float z) //camera's view direction -* UP (float x) (float y) (float z) //camera's up vector - -Objects are defined in the following fashion: - -* OBJECT (object ID) //object header -* (cube OR sphere OR mesh) //type of object, can be either "cube", "sphere", or - "mesh". Note that cubes and spheres are unit sized and centered at the - origin. -* material (material ID) //material to assign this object -* TRANS (float transx) (float transy) (float transz) //translation -* ROTAT (float rotationx) (float rotationy) (float rotationz) //rotation -* SCALE (float scalex) (float scaley) (float scalez) //scale - -Two examples are provided in the `scenes/` directory: a single emissive sphere, -and a simple cornell box made using cubes for walls and lights and a sphere in -the middle. - -## Third-Party Code Policy - -* Use of any third-party code must be approved by asking on our Google Group. -* If it is approved, all students are welcome to use it. Generally, we approve - use of third-party code that is not a core part of the project. For example, - for the path 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, at minimum, - result in you receiving an F for the semester. - -## README - -Please see: [**TIPS FOR WRITING AN AWESOME README**](https://github.com/pjcozzi/Articles/blob/master/CIS565/GitHubRepo/README.md) - -* Sell your project. -* Assume the reader has a little knowledge of path tracing - don't go into - detail explaining what it is. Focus on your project. -* Don't talk about it like it's an assignment - don't say what is and isn't - "extra" or "extra credit." Talk about what you accomplished. -* Use this to document what you've done. -* *DO NOT* leave the README to the last minute! It is a crucial part of the - project, and we will not be able to grade you without a good README. - -In addition: - -* This is a renderer, so include images that you've made! -* Be sure to back your claims for optimization with numbers and comparisons. -* If you reference any other material, please provide a link to it. -* You wil not be graded on how fast your path tracer runs, but getting close to - real-time is always nice! -* If you have a fast GPU renderer, it is very good to show case this with a - video to show interactivity. If you do so, please include a link! +![](img/final1000.png "Final image after 1000 iterations") +* This final image shows the refracting surfaces, specular surfaces, difuse, and mirrors. ### Analysis @@ -278,24 +35,14 @@ In addition: effects of stream compaction within a single iteration (i.e. the number of unterminated rays after each bounce) and evaluate the benefits you get from stream compaction. + ![](img/Proj3_chart1.png "Threads vs. Trace Depth") + ![](img/Proj3_chart2.png "Milliseconds vs. Trace Depth") + * As you can see from the above graphs, as the number of threads decreases, the amount of time spend in each path trace decreases. The green line shows when there is no stream compaction. It is clear why the time stays fairly consistant, becuase there is no change in the number of threads being used on the GPU. It starts off taking less time than the stream compaction because of the memory transfers that the stream compaction must do. However, in the end, the stream compaction saves enough time, by lowering the number of threads used in each path trace, that the overall time to run is lower when stream compaction is used. It is also clear from these graphs, that stream compaction allows for more of a speer-up in an open room. This is talked about more in the below paragraph. * Compare scenes which are open (like the given cornell box) and closed (i.e. no light can escape the scene). Again, compare the performance effects of stream compaction! Remember, stream compaction only affects rays which terminate, so what might you expect? + * Because rays terminate when they hit a light, or do not intersect with anything, in a closed room, there will not be any rays that do not intersect anything. Thus, less rays will be terminated. It is because of this that the stream compaction will not be used as much as compared to an open room. We would guess that stream compaction does not help speed up the running time in a closed room as much as it does in an open one. Test results helped to prove this answer. The average time that one iteration took in an open room with no stream compaction was 578.326 milliseconds. The average time for an open room with stream compaction was 467.913. This means that the average speed-up time that stream compaction gave in an open room was 110.413 milliseconds per iteration. Now we can compare this to the average time in a closed room. The average time for a closed room with no stream compaction was 757.264 milliseconds per iteration. For a closed room with stream compaction it was 664.501 milliseconds per iteration. This means that the average speed-up time for a closed room was only 92.763 milliseconds per iteration. This shows how stream compaction is more helpful in a open room compared to a closed room. -## Submit - -If you have modified any of the `CMakeLists.txt` files at all (aside from the -list of `SOURCE_FILES`), you must test that your project can build in Moore -100B/C. Beware of any build issues discussed on the Google Group. -1. Open a GitHub pull request so that we can see that you have finished. - The title should be "Submission: YOUR NAME". -2. Send an email to the TA (gmail: kainino1+cis565@) with: - * **Subject**: in the form of `[CIS565] Project N: PENNKEY`. - * Direct link to your pull request on GitHub. - * Estimate the amount of time you spent on the project. - * If there were any outstanding problems, or if you did any extra - work, *briefly* explain. - * Feedback on the project itself, if any. diff --git a/img/Closed box.png b/img/Closed box.png new file mode 100644 index 0000000..380e705 Binary files /dev/null and b/img/Closed box.png differ diff --git a/img/DOF1.png b/img/DOF1.png new file mode 100644 index 0000000..1c3164c Binary files /dev/null and b/img/DOF1.png differ diff --git a/img/DOF2.png b/img/DOF2.png new file mode 100644 index 0000000..1074195 Binary files /dev/null and b/img/DOF2.png differ diff --git a/img/DOF_1000.png b/img/DOF_1000.png new file mode 100644 index 0000000..838d55f Binary files /dev/null and b/img/DOF_1000.png differ diff --git a/img/DOF_500.png b/img/DOF_500.png new file mode 100644 index 0000000..1962d24 Binary files /dev/null and b/img/DOF_500.png differ diff --git a/img/Jittered with -.5, .5.png b/img/Jittered with -.5, .5.png new file mode 100644 index 0000000..3055fc4 Binary files /dev/null and b/img/Jittered with -.5, .5.png differ diff --git a/img/Jittered with 0, 1.png b/img/Jittered with 0, 1.png new file mode 100644 index 0000000..d55ffbc Binary files /dev/null and b/img/Jittered with 0, 1.png differ diff --git a/img/Proj3_chart1.png b/img/Proj3_chart1.png new file mode 100644 index 0000000..2c2fc37 Binary files /dev/null and b/img/Proj3_chart1.png differ diff --git a/img/Proj3_chart2.png b/img/Proj3_chart2.png new file mode 100644 index 0000000..85549cc Binary files /dev/null and b/img/Proj3_chart2.png differ diff --git a/img/cornell.2015-09-24_21-39-23z.168samp.png b/img/cornell.2015-09-24_21-39-23z.168samp.png new file mode 100644 index 0000000..4cdcf81 Binary files /dev/null and b/img/cornell.2015-09-24_21-39-23z.168samp.png differ diff --git a/img/cornell.2015-09-24_22-03-22z.145samp.png b/img/cornell.2015-09-24_22-03-22z.145samp.png new file mode 100644 index 0000000..dc16eef Binary files /dev/null and b/img/cornell.2015-09-24_22-03-22z.145samp.png differ diff --git a/img/cornell.2015-09-24_22-46-43z.148samp.png b/img/cornell.2015-09-24_22-46-43z.148samp.png new file mode 100644 index 0000000..10714ff Binary files /dev/null and b/img/cornell.2015-09-24_22-46-43z.148samp.png differ diff --git a/img/cornell.2015-09-25_22-03-49z.862samp.png b/img/cornell.2015-09-25_22-03-49z.862samp.png new file mode 100644 index 0000000..cb78a08 Binary files /dev/null and b/img/cornell.2015-09-25_22-03-49z.862samp.png differ diff --git a/img/final1.png b/img/final1.png new file mode 100644 index 0000000..0742a8e Binary files /dev/null and b/img/final1.png differ diff --git a/img/final1000.png b/img/final1000.png new file mode 100644 index 0000000..8e96691 Binary files /dev/null and b/img/final1000.png differ diff --git a/img/motion_blur1.png b/img/motion_blur1.png new file mode 100644 index 0000000..4980664 Binary files /dev/null and b/img/motion_blur1.png differ diff --git a/img/motion_blur_1000.png b/img/motion_blur_1000.png new file mode 100644 index 0000000..e4eb814 Binary files /dev/null and b/img/motion_blur_1000.png differ diff --git a/img/no jitter.png b/img/no jitter.png new file mode 100644 index 0000000..71aae43 Binary files /dev/null and b/img/no jitter.png differ diff --git a/img/old seeds.png b/img/old seeds.png new file mode 100644 index 0000000..1d2fc69 Binary files /dev/null and b/img/old seeds.png differ diff --git a/img/refraction w prob.png b/img/refraction w prob.png new file mode 100644 index 0000000..b585b04 Binary files /dev/null and b/img/refraction w prob.png differ diff --git a/img/refraction1.png b/img/refraction1.png new file mode 100644 index 0000000..bed7047 Binary files /dev/null and b/img/refraction1.png differ diff --git a/img/refraction2 large light.png b/img/refraction2 large light.png new file mode 100644 index 0000000..9dedb51 Binary files /dev/null and b/img/refraction2 large light.png differ diff --git a/img/refraction2.png b/img/refraction2.png new file mode 100644 index 0000000..024fef0 Binary files /dev/null and b/img/refraction2.png differ diff --git a/img/refraction3.png b/img/refraction3.png new file mode 100644 index 0000000..7340c58 Binary files /dev/null and b/img/refraction3.png differ diff --git a/img/refraction_correct1.png b/img/refraction_correct1.png new file mode 100644 index 0000000..64a3b4a Binary files /dev/null and b/img/refraction_correct1.png differ diff --git a/img/refraction_correct2(water).png b/img/refraction_correct2(water).png new file mode 100644 index 0000000..1c5e067 Binary files /dev/null and b/img/refraction_correct2(water).png differ diff --git a/img/refraction_correct3(glass).png b/img/refraction_correct3(glass).png new file mode 100644 index 0000000..2a949b0 Binary files /dev/null and b/img/refraction_correct3(glass).png differ diff --git a/img/refraction_correctish.png b/img/refraction_correctish.png new file mode 100644 index 0000000..2d2bc7f Binary files /dev/null and b/img/refraction_correctish.png differ diff --git a/img/refraction_w_light_box.png b/img/refraction_w_light_box.png new file mode 100644 index 0000000..f809b88 Binary files /dev/null and b/img/refraction_w_light_box.png differ diff --git a/img/refraction_w_light_sphere.png b/img/refraction_w_light_sphere.png new file mode 100644 index 0000000..bcb4be6 Binary files /dev/null and b/img/refraction_w_light_sphere.png differ diff --git a/img/specular1.png b/img/specular1.png new file mode 100644 index 0000000..05112b2 Binary files /dev/null and b/img/specular1.png differ diff --git a/img/specular2.png b/img/specular2.png new file mode 100644 index 0000000..2a20999 Binary files /dev/null and b/img/specular2.png differ diff --git a/img/specular3.png b/img/specular3.png new file mode 100644 index 0000000..b100d4d Binary files /dev/null and b/img/specular3.png differ diff --git a/img/unterminated rays to black.png b/img/unterminated rays to black.png new file mode 100644 index 0000000..7ed589a Binary files /dev/null and b/img/unterminated rays to black.png differ diff --git a/scenes/cornell.txt b/scenes/cornell.txt index a18639e..5b0dd3c 100644 --- a/scenes/cornell.txt +++ b/scenes/cornell.txt @@ -6,7 +6,7 @@ SPECRGB 0 0 0 REFL 0 REFR 0 REFRIOR 0 -EMITTANCE 5 +EMITTANCE 2 // Diffuse white MATERIAL 1 @@ -38,6 +38,36 @@ REFR 0 REFRIOR 0 EMITTANCE 0 +// Mirror +MATERIAL 4 +RGB 1 1 1 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 1 +REFRIOR 1.52 +EMITTANCE 0 + +// Mirror +MATERIAL 5 +RGB 1 1 1 +SPECEX 0 +SPECRGB 0 1 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Mirror +MATERIAL 6 +RGB 1 1 1 +SPECEX 0 +SPECRGB 1 1 1 +REFL 1 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + // Camera CAMERA RES 800 800 @@ -56,7 +86,7 @@ cube material 0 TRANS 0 10 0 ROTAT 0 0 0 -SCALE 3 .3 3 +SCALE 4 .3 4 // Floor OBJECT 1 @@ -65,14 +95,16 @@ material 1 TRANS 0 0 0 ROTAT 0 0 0 SCALE 10 .01 10 +MOVE 0 0 0 // Ceiling OBJECT 2 cube -material 1 +material 0 TRANS 0 10 0 ROTAT 0 0 90 SCALE .01 10 10 +MOVE 0 0 0 // Back wall OBJECT 3 @@ -81,6 +113,7 @@ material 1 TRANS 0 5 -5 ROTAT 0 90 0 SCALE .01 10 10 +MOVE 0 0 0 // Left wall OBJECT 4 @@ -89,6 +122,7 @@ material 2 TRANS -5 5 0 ROTAT 0 0 0 SCALE .01 10 10 +MOVE 0 0 0 // Right wall OBJECT 5 @@ -97,11 +131,35 @@ material 3 TRANS 5 5 0 ROTAT 0 0 0 SCALE .01 10 10 +MOVE 0 0 0 // Sphere OBJECT 6 sphere -material 1 +material 6 TRANS -1 4 -1 ROTAT 0 0 0 +SCALE 2 2 2 +MOVE 0 0 0 + +// Sphere +OBJECT 7 +cube +material 4 +TRANS 2 1 -1 +ROTAT 0 45 0 +SCALE 2 2 2 +MOVE 0 0 0 + +// Sphere +OBJECT 8 +sphere +material 5 +TRANS -3 1 -1 +ROTAT 0 45 0 SCALE 3 3 3 +MOVE 0 0 0 + + + + diff --git a/scenes/cornellTest.txt b/scenes/cornellTest.txt new file mode 100644 index 0000000..3181233 --- /dev/null +++ b/scenes/cornellTest.txt @@ -0,0 +1,149 @@ +// Emissive material (light) +MATERIAL 0 +RGB 1 1 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 2 + +// Diffuse white +MATERIAL 1 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Diffuse red +MATERIAL 2 +RGB .85 .35 .35 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Diffuse green +MATERIAL 3 +RGB .35 .85 .35 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Mirror +MATERIAL 4 +RGB 1 1 1 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 1 +REFRIOR 1.52 +EMITTANCE 0 + +// Mirror +MATERIAL 5 +RGB 0 0 1 +SPECEX 1 +SPECRGB 1 1 1 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Mirror +MATERIAL 6 +RGB 0 1 1 +SPECEX 1 +SPECRGB 1 1 1 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Camera +CAMERA +RES 800 800 +FOVY 45 +ITERATIONS 5000 +DEPTH 8 +FILE cornell +EYE 0.0 5 10.5 +VIEW 0 0 -1 +UP 0 1 0 + + +// Ceiling light +OBJECT 0 +cube +material 0 +TRANS 0 10 0 +ROTAT 0 0 0 +SCALE 4 .3 4 + +// Floor +OBJECT 1 +cube +material 1 +TRANS 0 0 0 +ROTAT 0 0 0 +SCALE 10 .01 10 +MOVE 0 0 0 + +// Ceiling +OBJECT 2 +cube +material 0 +TRANS 0 10 0 +ROTAT 0 0 90 +SCALE .01 10 10 +MOVE 0 0 0 + +// Back wall +OBJECT 3 +cube +material 1 +TRANS 0 5 -5 +ROTAT 0 90 0 +SCALE .01 10 10 +MOVE 0 0 0 + +// Left wall +OBJECT 4 +cube +material 2 +TRANS -5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 +MOVE 0 0 0 + +// Right wall +OBJECT 5 +cube +material 3 +TRANS 5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 +MOVE 0 0 0 + +// Sphere +OBJECT 6 +sphere +material 6 +TRANS -1 4 -1 +ROTAT 0 0 0 +SCALE 3 3 3 +MOVE 0 0 0 + + + + + + diff --git a/scenes/sphere.txt b/scenes/sphere.txt index c70d546..3a57568 100644 --- a/scenes/sphere.txt +++ b/scenes/sphere.txt @@ -10,7 +10,7 @@ EMITTANCE 5 // Camera CAMERA -RES 800 800 +RES 10 10 FOVY 45 ITERATIONS 5000 DEPTH 8 @@ -19,10 +19,12 @@ EYE 0.0 5 10.5 VIEW 0 0 -1 UP 0 1 0 -// Sphere +// sphere OBJECT 0 sphere material 0 TRANS 0 0 0 ROTAT 0 0 0 SCALE 3 3 3 +MOVE 0 2 0 + diff --git a/src/interactions.h b/src/interactions.h index 9386a09..2057de5 100644 --- a/src/interactions.h +++ b/src/interactions.h @@ -71,8 +71,76 @@ void scatterRay( glm::vec3 intersect, glm::vec3 normal, const Material &m, + bool out, thrust::default_random_engine &rng) { // TODO: implement this. // A basic implementation of pure-diffuse shading will just call the // calculateRandomDirectionInHemisphere defined above. + + + if (m.hasRefractive == 1) { + thrust::uniform_real_distribution probDistrib(0.0f, 1.0f); + float prob = probDistrib(rng); + float angle; + glm::vec3 refractionPoint; + + + float R_0; + if (out) { + R_0 = ((1.0f - m.indexOfRefraction) / (1.0f + m.indexOfRefraction))*((1.0f - m.indexOfRefraction) / (1.0f + m.indexOfRefraction)); + float n_12 = glm::pow(1.0f / m.indexOfRefraction, 2); + angle = 1.0f - n_12 * (1.0f - glm::pow(glm::dot(normal, ray.direction), 2)); + } + else { + R_0 = ((m.indexOfRefraction - 1.0f) / (m.indexOfRefraction + 1.0f)) * ((m.indexOfRefraction - 1.0f) / (m.indexOfRefraction + 1.0f)); + angle = 1.0f - glm::pow(m.indexOfRefraction, 2) * (1.0f - glm::pow(glm::dot(normal, ray.direction), 2)); + } + float reflCoeff = R_0 + (1.0f - R_0) * glm::pow(1.0f - glm::dot(normal, -ray.direction), 5); + if ((1.0f - reflCoeff) > prob && angle > 0.0f) { + if (out == true){ //if ray is coming from air to geom + refractionPoint = glm::refract(ray.direction, normal, 1.0f / m.indexOfRefraction); + ray.out = false; + } + else{ //if ray is coming out of geom into air + refractionPoint = glm::refract(ray.direction, normal, m.indexOfRefraction); + ray.out = true; + } + ray.direction = refractionPoint; + ray.origin = intersect + glm::vec3(0.01f, 0.01f, 0.01f)*(glm::normalize(ray.direction)); + color *= m.color;// *(1.0f / reflCoeff); + + } + else { + ray.direction = ray.direction - 2.0f*normal*(glm::dot(ray.direction, normal)); + ray.origin = intersect + glm::vec3(0.01f, 0.01f, 0.01f)*(glm::normalize(ray.direction)); + color *= m.color;// *(1.0f / reflCoeff); + } + + } + else if (m.hasReflective == 1) { + ray.direction = ray.direction - 2.0f*normal*(glm::dot(ray.direction, normal)); + ray.origin = intersect + glm::vec3(0.01f, 0.01f, 0.01f)*(glm::normalize(ray.direction)); + color *= m.color; + } + else if (m.specular.exponent > 0) { + thrust::uniform_real_distribution probDistrib(0.0f, 1.0f); + float prob = probDistrib(rng); + if (.1f > prob) { + ray.direction = ray.direction - 2.0f*normal*(glm::dot(ray.direction, normal)); + ray.origin = intersect + glm::vec3(0.01f, 0.01f, 0.01f)*(glm::normalize(ray.direction)); + color *= m.specular.color; + } + else { + ray.direction = calculateRandomDirectionInHemisphere(normal, rng); + ray.origin = intersect + glm::vec3(0.01f, 0.01f, 0.01f)*(glm::normalize(ray.direction)); + color *= m.color; + } + } + //DIFFUSE + else { + ray.direction = calculateRandomDirectionInHemisphere(normal, rng); + ray.origin = intersect + glm::vec3(0.01f, 0.01f, 0.01f)*(glm::normalize(ray.direction)); + color *= m.color; + } + } diff --git a/src/pathtrace.cu b/src/pathtrace.cu index a1d9e24..bdf6d7f 100644 --- a/src/pathtrace.cu +++ b/src/pathtrace.cu @@ -1,9 +1,12 @@ #include #include +#include #include #include #include #include +#include +#include #include "sceneStructs.h" #include "scene.h" @@ -16,7 +19,21 @@ #define FILENAME (strrchr(__FILE__, '/') ? strrchr(__FILE__, '/') + 1 : __FILE__) #define checkCUDAError(msg) checkCUDAErrorFn(msg, FILENAME, __LINE__) + +inline int ilog2(int x) { + int lg = 0; + while (x >>= 1) { + ++lg; + } + return lg; +} + +inline int ilog2ceil(int x) { + return ilog2(x - 1) + 1; +} + void checkCUDAErrorFn(const char *msg, const char *file, int line) { + cudaDeviceSynchronize(); cudaError_t err = cudaGetLastError(); if (cudaSuccess == err) { return; @@ -61,6 +78,13 @@ __global__ void sendImageToPBO(uchar4* pbo, glm::ivec2 resolution, static Scene *hst_scene = NULL; static glm::vec3 *dev_image = NULL; +static Geom *dev_geoms = NULL; +static Material *dev_mats = NULL; +static Ray *dev_rayArray = NULL; + +int* dev_bools; +int* dev_indices; + // TODO: static variables for device memory, scene/camera info, etc // ... @@ -68,9 +92,20 @@ void pathtraceInit(Scene *scene) { hst_scene = scene; const Camera &cam = hst_scene->state.camera; const int pixelcount = cam.resolution.x * cam.resolution.y; + const Geom *geoms = &(hst_scene->geoms)[0]; + const Material *mats = &(hst_scene->materials)[0]; + cudaMalloc(&dev_geoms, pixelcount * sizeof(Geom)); + cudaMalloc(&dev_mats, pixelcount * sizeof(Material)); + cudaMalloc(&dev_rayArray, pixelcount * sizeof(Ray)); cudaMalloc(&dev_image, pixelcount * sizeof(glm::vec3)); cudaMemset(dev_image, 0, pixelcount * sizeof(glm::vec3)); + cudaMemset(dev_rayArray, 0, pixelcount * sizeof(Ray)); + + + cudaMemcpy(dev_mats, mats, hst_scene->materials.size() * sizeof(Material), cudaMemcpyHostToDevice); + + // TODO: initialize the above static variables added above checkCUDAError("pathtraceInit"); @@ -79,7 +114,9 @@ void pathtraceInit(Scene *scene) { void pathtraceFree() { cudaFree(dev_image); // no-op if dev_image is null // TODO: clean up the above static variables - + cudaFree(dev_geoms); + cudaFree(dev_mats); + cudaFree(dev_rayArray); checkCUDAError("pathtraceFree"); } @@ -108,21 +145,421 @@ __global__ void generateNoiseDeleteMe(Camera cam, int iter, glm::vec3 *image) { } } +//Create ray to be shot at a pixel in the image +__global__ void kernRayGenerate(Camera cam, Ray *ray, int iter, bool dof){ + int x = (blockIdx.x * blockDim.x) + threadIdx.x; + int y = (blockIdx.y * blockDim.y) + threadIdx.y; + int index = x + (y*cam.resolution.x); + thrust::default_random_engine rng = makeSeededRandomEngine(iter, index, 0); + thrust::uniform_real_distribution unitDistrib(-.5f, .5f); + thrust::uniform_real_distribution dofDistrib(-1.0f, 1.0f); + //Calculate camera's world position + if (x < cam.resolution.x && y < cam.resolution.y) { + glm::vec3 A = glm::cross(cam.view, cam.up); + glm::vec3 B = glm::cross(A, cam.view); + glm::vec3 M = cam.position + cam.view; + float lenC = glm::length(cam.view); + float lenA = glm::length(A); + float lenB = glm::length(B); + float tantheta = (float)cam.resolution.x; + tantheta /= (float)cam.resolution.y; + tantheta *= tan((float)glm::radians(cam.fov[1])); + + glm::vec3 H = (A*lenC*tantheta) / lenA; + glm::vec3 V = (B*lenC*tan((float)glm::radians(cam.fov[1]))) / lenB; + + //Create ray with direction and origin + //Jitter rays with uniform distribution + //printf("%f ", unitDistrib(rng)); + float sx = ((float)x + unitDistrib(rng)) / ((float)cam.resolution.x - 1.0f); + float sy = ((float)y + unitDistrib(rng)) / ((float)cam.resolution.y - 1.0f); + //Get world coordinates of pixel + glm::vec3 WC = M - (2.0f*sx - 1.0f)*H - (2.0f*sy - 1.0f)*V; + //Get direction of ray + glm::vec3 dir = glm::normalize(WC - cam.position); + + ray[index].origin = cam.position; + ray[index].direction = dir; + ray[index].color = glm::vec3(1.0, 1.0, 1.0); + ray[index].index = index; + ray[index].terminated = false; + ray[index].out = true; + if (dof == true) { + glm::vec3 apOff = glm::vec3(dofDistrib(rng), dofDistrib(rng), 0.0f); + glm::vec3 new_E = cam.position + apOff; + float focal = 11.5866f; //glm::length(glm::vec3(-2.0f, 5.0f,2.0f) - new_E); + dir *= focal; + dir -= apOff; + dir = glm::normalize(dir); + ray[index].origin = new_E; + ray[index].direction = dir; + } + } + +} + +//Helper function to get random point on cubic light +__device__ glm::vec3 getRandomPointOnCube(Geom node, int iter, int index) { + // TODO: get the dimensions of the transformed cube in world space + glm::vec3 dim(0.0f, 0.0f, 0.0f); + dim = node.scale; + + // Get surface area of the cube + float side1 = dim[0] * dim[1]; // x-y + float side2 = dim[1] * dim[2]; // y-z + float side3 = dim[0] * dim[2]; // x-z + float totalArea = 2.0f * (side1 + side2 + side3); + + thrust::default_random_engine rng = makeSeededRandomEngine(iter, index, 0); + thrust::uniform_real_distribution unitDistrib(-.5f, .5f); + thrust::uniform_real_distribution dofDistrib(0.0f, 1.0f); + + // pick random face weighted by surface area + float r = floor(dofDistrib(rng)); + // pick 2 random components for the point in the range (-0.5, 0.5) + float c1 = unitDistrib(rng); + float c2 = unitDistrib(rng); + + glm::vec3 point; + if (r < side1 / totalArea) { + // x-y front + point = glm::vec3(c1, c2, 0.5f); + } else if (r < (side1 * 2) / totalArea) { + // x-y back + point = glm::vec3(c1, c2, -0.5f); + } else if (r < (side1 * 2 + side2) / totalArea) { + // y-z front + point = glm::vec3(0.5f, c1, c2); + } else if (r < (side1 * 2 + side2 * 2) / totalArea) { + // y-z back + point = glm::vec3(-0.5f, c1, c2); + } else if (r < (side1 * 2 + side2 * 2 + side3) / totalArea) { + // x-z front + point = glm::vec3(c1, 0.5f, c2); + } else { + // x-z back + point = glm::vec3(c1, -0.5f, c2); + } + + // TODO: transform point to world space + glm::mat4 T(1.0f); + T = glm::translate(T, node.translation); + + if (node.rotation[0] != 0){ + T = glm::rotate(T, node.rotation[0]*(PI/180.0f), glm::vec3(1,0,0)); + } + if (node.rotation[1] != 0){ + T = glm::rotate(T, node.rotation[1]*(PI/180.0f), glm::vec3(0,1,0)); + } + if (node.rotation[2] != 0){ + T = glm::rotate(T, node.rotation[2]*(PI/180.0f), glm::vec3(0,0,1)); + } + + //T = glm::scale(T, node.scale); + glm::vec4 newPoint = T*glm::vec4(point, 1.0f); + point = glm::vec3(newPoint[0], newPoint[1], newPoint[2]); + return point; +} + +//Helper function to get random point on spherical light +/*__device__ glm::vec3 getRandomPointOnSphere(Geom node, int iter, int index) { + // generate u, v, in the range (0, 1) + float u = static_cast (rand()) / static_cast (RAND_MAX); + float v = static_cast (rand()) / static_cast (RAND_MAX); + + float theta = 2.0f * PI * u; + float phi = acos(2.0f * v - 1.0f); + + // find x, y, z coordinates assuming unit sphere in object space + glm::vec3 point; + point[0] = sin(phi) * cos(theta); + point[1] = sin(phi) * sin(theta); + point[2] = cos(phi); + + // TODO: transform point to world space + glm::mat4 T(1.0f); + T = glm::translate(T, node.translation); + + if (node.rotation[0] != 0){ + T = glm::rotate(T, node.rotation[0]*(PI/180.0f), glm::vec3(1,0,0)); + } + if (node.rotation[1] != 0){ + T = glm::rotate(T, node.rotation[1]*(PI/180.0f), glm::vec3(0,1,0)); + } + if (node.rotation[2] != 0){ + T = glm::rotate(T, node.rotation[2]*(PI/180.0f), glm::vec3(0,0,1)); + } + + glm::vec4 newPoint = T*glm::vec4(point, 1.0f); + point = glm::vec3(newPoint[0], newPoint[1], newPoint[2]); + return point; +}*/ +//Helper function to find closest intersection +__device__ float closestIntersection(Ray ray, const Geom* geoms, glm::vec3 &intersectionPoint, glm::vec3 &normal, bool &outside, int &objIndex, const int numGeoms){ + glm::vec3 interPoint; + glm::vec3 norm; + bool out; + float t = -1; + float dist; + for (int i = 0; i < numGeoms; i++) { + if (geoms[i].type == CUBE) { + dist = boxIntersectionTest(geoms[i], ray, interPoint, norm, out); + } + else if (geoms[i].type == SPHERE) { + dist = sphereIntersectionTest(geoms[i], ray, interPoint, norm, out); + } + if ((dist != -1 && dist < t) || t == -1) { + t = dist; + intersectionPoint = interPoint; + normal = norm; + outside = out; + objIndex = i; + + } + } + return t; + +} + +//Function to find next ray +__global__ void kernPathTracer(Camera cam, Ray* rayArray, const Geom* geoms, const Material* mats, const int numGeoms, const int numMats, glm::vec3* dev_image, int iter, int depth, int traceDepth, bool m_blur, int size){ + //int x = (blockIdx.x * blockDim.x) + threadIdx.x; + //int y = (blockIdx.y * blockDim.y) + threadIdx.y; + //int index = x + (y * cam.resolution.x); + int index = (blockIdx.x * blockDim.x) + threadIdx.x; + int imageSize = (cam.resolution.x * cam.resolution.y); + //find closest intersection + /*if (rayArray[index].terminated == true && index < size) { + dev_image[rayArray[index].index] = glm::vec3(1.0, 0.0, 0.0); + printf("in here: %i, %i, %i \n", size, index, rayArray[index].index); + return; + } + else { + //printf("is fine: %i, %i, %i \n", depth, index, rayArray[index].index); + }*/ + + if (index < size && rayArray[index].terminated == false) {//rayArray[index].index < imageSize && index < size && rayArray[index].terminated == false) {//x < cam.resolution.x && y < cam.resolution.y && rayArray[index].terminated == false) { + thrust::default_random_engine rng = makeSeededRandomEngine(iter, index, depth); + glm::vec3 interPoint; + glm::vec3 norm; + bool out; + int objIndex; + if (depth == traceDepth) { + dev_image[rayArray[index].index] == glm::vec3(0.0f, 0.0f, 0.0f); + return; + /*for (int i = 0; i < numGeoms; i++) { + if (mats[geoms[i].materialid].emittance > 0 && mats[geoms[rayArray[index].geomid].materialid].emittance == 0 && mats[geoms[rayArray[index].geomid].materialid].hasReflective == 0 && mats[geoms[rayArray[index].geomid].materialid].hasRefractive == 0) { + glm::vec3 new_pt = getRandomPointOnCube(geoms[i], iter, index); + rayArray[index].direction = rayArray[index].origin + glm::normalize(new_pt - rayArray[index].origin); + float t = closestIntersection(rayArray[index], geoms, interPoint, norm, out, objIndex, numGeoms); + if (objIndex == i) { + printf("hit light in direct"); + rayArray[index].color *= mats[geoms[i].materialid].emittance*mats[geoms[objIndex].materialid].color; + dev_image[index] += rayArray[index].color; + } + } + }*/ + + } + //Geom* m_blur_geoms = new Geom[numGeoms]; + float t; + + if (m_blur) { + /*for (int i = 0; i < numGeoms; i++) { + m_blur_geoms[i] = geoms[i]; + m_blur_geoms[i].translation.x += m_blur_geoms[i].move.x*rayArray[index].time; + m_blur_geoms[i].translation.y += m_blur_geoms[i].move.y*rayArray[index].time; + m_blur_geoms[i].translation.z += m_blur_geoms[i].move.z*rayArray[index].time; + }*/ + t = closestIntersection(rayArray[index], geoms, interPoint, norm, out, objIndex, numGeoms); + } + else { + t = closestIntersection(rayArray[index], geoms, interPoint, norm, out, objIndex, numGeoms); + } + rayArray[index].geomid = objIndex; + //get direction of next ray and compute new color + if (t >= 0.0f) { + if (mats[geoms[objIndex].materialid].emittance >= 1) { + rayArray[index].color *= mats[geoms[objIndex].materialid].emittance*mats[geoms[objIndex].materialid].color; + dev_image[rayArray[index].index] += rayArray[index].color; + rayArray[index].terminated = true; + } + else { + scatterRay(rayArray[index], rayArray[index].color, interPoint, norm, mats[geoms[objIndex].materialid], out, rng); + } + } + else { + //dev_image[index] *= glm::vec3(0.0f, 0.0f, 0.0f); //rayArray[index].color; + rayArray[index].terminated = true; + } + } + +} + + +__global__ void kernCombine(int *maxArray, int *newData, int n) { + int index = threadIdx.x + (blockIdx.x * blockDim.x); + + if (index < n) { + //printf("index: %i max: %i \n", g_idata[index], maxArray[blockIdx.x]); + newData[index] = newData[index] + maxArray[blockIdx.x]; + + } +} + +__global__ void kernScan(int *maxArray, int *g_idata, int n) { + //printf("IN THIN FUNCTION AELFHGGGGHF"); + extern __shared__ int temp[]; + //printf("blockId: %i", blockDim.x); + int thid = threadIdx.x + (blockIdx.x * blockDim.x); + int t = threadIdx.x; + int offset = 1; + temp[2*t] = g_idata[2*thid]; + temp[2*t+1] = g_idata[2*thid+1]; + for (int d = (2*blockDim.x)>>1; d > 0; d >>=1) { + __syncthreads(); + if (t < d) { + int ai = offset*(2*t+1)-1; + int bi = offset*(2*t+2)-1; + temp[bi] += temp[ai]; + } + offset *= 2; + } + if (t == 0) { + temp[2*blockDim.x-1] = 0; + } + + for (int d = 1; d < (2*blockDim.x); d*=2) { + offset >>= 1; + __syncthreads(); + if (t < d) { + int ai = offset *(2*t+1)-1; + int bi = offset *(2*t+2)-1; + float t2 = temp[ai]; + temp[ai] = temp[bi]; + temp[bi] += t2; + } + } + __syncthreads(); + if (t == (blockDim.x - 1)) { + maxArray[blockIdx.x] = temp[2*t+1] + g_idata[2*thid+1]; + } + g_idata[2*thid] = temp[2*t]; + //printf("(%i, %i) \n", thid, g_idata[2*thid]); + g_idata[2*thid+1] = temp[2*t+1]; + +} + +void scan(int n, const int *idata, int *odata, int newSize) { + int blockSize = 32; + int numBlocks = ceil((float)n / (float)blockSize); + int powTwo = 1<>>(maxArray, newArray, n); + checkCUDAError("pathtrace"); + cudaDeviceSynchronize(); + cudaMemcpy(odata, newArray, n*sizeof(int), cudaMemcpyDeviceToHost); + + int maxSize = ((powTwo/2) + blockSize - 1) / blockSize; + if (maxSize != 1) { + int* hst_maxArray = new int[maxSize]; + int* scanMax = new int[maxSize]; + int* dev_scanMax; + + cudaMalloc((void**)&dev_scanMax, maxSize*sizeof(int)); + + cudaMemcpy(hst_maxArray, maxArray, maxSize*sizeof(int), cudaMemcpyDeviceToHost); + //printf("%i ", hst_maxArray[maxSize - 1]); + + scan(maxSize, hst_maxArray, scanMax, randomNum); + //cudaMemcpy(odata, g_idata, n*sizeof(int), cudaMemcpyDeviceToHost); + + cudaMemcpy(dev_scanMax, scanMax, maxSize*sizeof(int), cudaMemcpyHostToDevice); + //kernCombine<<>>(dev_scanMax, g_idata, n); + //cudaMemcpy(odata, g_idata, n*sizeof(int), cudaMemcpyDeviceToHost); + cudaMemcpy(newArray, odata, n*sizeof(int), cudaMemcpyHostToDevice); + /*for (int i = 0; i < maxSize; i++) { + for (int j = blockSize*2*i; j < blockSize*2*(i+1)) { + odata[] + } + }*/ + kernCombine<<>>(dev_scanMax, newArray, n); + //checkCUDAError("pathtrace"); + cudaMemcpy(odata, newArray, n*sizeof(int), cudaMemcpyDeviceToHost); + newSize = hst_maxArray[maxSize - 1]; + cudaFree(dev_scanMax); + } + + + + //printf(" don with function "); + cudaFree(maxArray); + cudaFree(g_idata); + cudaFree(newArray); + //checkCUDAError("pathtrace"); +} + +__global__ void kernScatter(int n, Ray *odata, + const Ray *idata, const int *bools, const int *indices) { + int thrId = threadIdx.x + (blockIdx.x * blockDim.x); + if (thrId < n) { + //printf("BOOLS: %i \n", bools[thrId]); + if (bools[thrId] == 1) { + //printf("old index: %i , new index: %i \n", thrId, indices[thrId]); + + odata[indices[thrId]] = idata[thrId]; + //printf("old ray index: %i, new ray index: %i \n", idata[thrId].index, odata[indices[thrId]].index); + } + } +} + /** * Wrapper for the __global__ call that sets up the kernel calls and does a ton * of memory management */ void pathtrace(uchar4 *pbo, int frame, int iter) { + //printf("iter: %i \n", iter); const int traceDepth = hst_scene->state.traceDepth; const Camera &cam = hst_scene->state.camera; const int pixelcount = cam.resolution.x * cam.resolution.y; - + const Geom *geoms = &(hst_scene->geoms)[0]; + Geom *m_blur_geoms = &(hst_scene->geoms)[0]; + int numGeoms = hst_scene->geoms.size(); + int numMats = hst_scene->materials.size(); + Ray *rayArray = new Ray[pixelcount]; + int max_iter = 1000; //hst_scene->state.iterations; const int blockSideLength = 8; const dim3 blockSize(blockSideLength, blockSideLength); + checkCUDAError("pathtrace"); const dim3 blocksPerGrid( - (cam.resolution.x + blockSize.x - 1) / blockSize.x, - (cam.resolution.y + blockSize.y - 1) / blockSize.y); - + (cam.resolution.x + blockSize.x - 1) / blockSize.x, + (cam.resolution.y + blockSize.y - 1) / blockSize.y); + //const dim3 blocksPerGrid((pixelcount + blockSize - 1) / blockSize); /////////////////////////////////////////////////////////////////////////// // Recap: @@ -148,18 +585,157 @@ void pathtrace(uchar4 *pbo, int frame, int iter) { // * Finally, handle all of the paths that still haven't terminated. // (Easy way is to make them black or background-colored.) + // TODO: perform one iteration of path tracing + bool dof = true; + bool m_blur = false; + bool streamCompaction = true; + int size = pixelcount; + + if (m_blur && iter < max_iter) { + for (int i = 0; i < numGeoms; i++) { - generateNoiseDeleteMe<<>>(cam, iter, dev_image); + m_blur_geoms[i] = geoms[i]; + m_blur_geoms[i].translation.x += geoms[i].move.x / (float)max_iter; + m_blur_geoms[i].translation.y += geoms[i].move.y / (float)max_iter; + m_blur_geoms[i].translation.z += geoms[i].move.z / (float)max_iter; + m_blur_geoms[i].transform = utilityCore::buildTransformationMatrix(m_blur_geoms[i].translation, m_blur_geoms[i].rotation, m_blur_geoms[i].scale); + m_blur_geoms[i].inverseTransform = glm::inverse(m_blur_geoms[i].transform); + m_blur_geoms[i].invTranspose = glm::inverseTranspose(m_blur_geoms[i].transform); + //printf("(%f, %f, %f)", m_blur_geoms[i].translation.x, m_blur_geoms[i].translation.y, m_blur_geoms[i].translation.z); + } + cudaMemcpy(dev_geoms, m_blur_geoms, hst_scene->geoms.size() * sizeof(Geom), cudaMemcpyHostToDevice); + checkCUDAError("pathtrace"); + } + else { + cudaMemcpy(dev_geoms, geoms, hst_scene->geoms.size() * sizeof(Geom), cudaMemcpyHostToDevice); + } + checkCUDAError("pathtrace"); + int newSize = pixelcount; + //printf("Heyo starting over \n newSize: %i \n", newSize); + Ray* dev_rayShort; + cudaMalloc((void**)&dev_rayShort, pixelcount * sizeof(Ray)); + checkCUDAError("pathtrace"); + cudaMalloc((void**)&dev_rayArray, pixelcount * sizeof(Ray)); + checkCUDAError("pathtrace"); + kernRayGenerate<<>>(cam, dev_rayArray, iter, dof); + checkCUDAError("pathtrace"); + //cuda events + cudaEvent_t start, stop; + cudaEventCreate(&start); + cudaEventCreate(&stop); + cudaEventRecord(start); + checkCUDAError("pathtrace"); + for (int i = 0; i < traceDepth + 1; i++) { + + int* hst_bool = new int[newSize]; + int* dev_bool; + int* dev_indices; + int* hst_indices = new int[newSize]; + int blockSizeNew = 64; + dim3 blocksGridNew((newSize + blockSizeNew - 1) / blockSizeNew); + checkCUDAError("pathtrace"); + //printf("BlocksPerGrid: %i \n", (newSize + blockSizeNew - 1) / blockSizeNew); + if (newSize == 0) { + cudaEventRecord(stop); + break; + } + //printf("size: %i \n", newSize); + kernPathTracer<<>>(cam, dev_rayArray, dev_geoms, dev_mats, numGeoms, numMats, dev_image, iter, i, traceDepth, m_blur, newSize); + checkCUDAError("pathtrace"); + cudaMemcpy(rayArray, dev_rayArray, pixelcount*sizeof(Ray), cudaMemcpyDeviceToHost); + checkCUDAError("pathtrace"); + if (streamCompaction) { + /* + for (int m = 0; m < newSize; m++) { + float ran = rand()%100; + if (ran < 50) { + hst_bool[m] = 0; + //printf("%i: %i \n", m, hst_bool[m]); + } + else { + hst_bool[m] = 1; + //printf("%i: %i \n", m, hst_bool[m]); + } + } + if (newSize > 1) { + scan(newSize, hst_bool, hst_indices, newSize); + newSize = hst_bool[newSize - 1] + hst_indices[newSize - 1]; + }*/ + + for (int m = 0; m < newSize; m++) { + + if (rayArray[m].terminated) { + hst_bool[m] = 0; + + //printf("%i: %i \n", m, hst_bool[m]); + + } else { + hst_bool[m] = 1; + + //printf("%i: %i \n", m, hst_bool[m]); + + //printf("%i: %i \n", m, hst_bool[m]); + } + } + int oldSize = newSize; + if (newSize > 1) { + //printf("old Size: %i \n", oldSize ); + + scan(oldSize, hst_bool, hst_indices, newSize); + + newSize = hst_bool[oldSize - 1] + hst_indices[oldSize - 1]; + //printf("new Size: %i \n", newSize ); + + cudaMalloc((void**)&dev_bool, oldSize * sizeof(int)); + cudaMalloc((void**)&dev_indices, oldSize * sizeof(int)); + + + cudaMemcpy(dev_bool, hst_bool, oldSize*sizeof(int), cudaMemcpyHostToDevice); + cudaMemcpy(dev_indices, hst_indices, oldSize*sizeof(int), cudaMemcpyHostToDevice); + checkCUDAError("pathtrace"); + //printf("%i, %i, %i \n", oldSize, blocksGridNew, blockSizeNew); + kernScatter<<>>(oldSize, dev_rayShort, dev_rayArray, dev_bool, dev_indices); + checkCUDAError("pathtrace"); + //printf("newSize: %i \n", newSize); + //cudaMemcpy(rayArray, dev_rayShort, newSize*sizeof(Ray), cudaMemcpyDeviceToHost); + + cudaMemcpy(dev_rayArray, dev_rayShort, newSize*sizeof(Ray), cudaMemcpyDeviceToDevice); + checkCUDAError("pathtrace"); + cudaFree(dev_bool); + cudaFree(dev_indices); + + } + + int k; + //std::cin >> k; + } + + + } + cudaEventRecord(stop); + cudaEventSynchronize(stop); + float milliseconds = 0; + cudaEventElapsedTime(&milliseconds, start, stop); + //printf("time per iteration: %f \n", milliseconds); + + + checkCUDAError("pathtrace"); + + cudaMemcpy(rayArray, dev_rayArray, pixelcount*sizeof(Ray), cudaMemcpyDeviceToHost); + checkCUDAError("pathtrace"); + //generateNoiseDeleteMe<<>>(cam, iter, dev_image); + /////////////////////////////////////////////////////////////////////////// // Send results to OpenGL buffer for rendering sendImageToPBO<<>>(pbo, cam.resolution, iter, dev_image); - + checkCUDAError("pathtrace"); // Retrieve image from GPU cudaMemcpy(hst_scene->state.image.data(), dev_image, pixelcount * sizeof(glm::vec3), cudaMemcpyDeviceToHost); - + cudaFree(dev_rayShort); + cudaFree(dev_rayArray); checkCUDAError("pathtrace"); } diff --git a/src/scene.cpp b/src/scene.cpp index 5804ce3..594124e 100644 --- a/src/scene.cpp +++ b/src/scene.cpp @@ -74,7 +74,9 @@ int Scene::loadGeom(string objectid) { newGeom.rotation = glm::vec3(atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str())); } else if (strcmp(tokens[0].c_str(), "SCALE") == 0) { newGeom.scale = glm::vec3(atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str())); - } + } else if (strcmp(tokens[0].c_str(), "MOVE") == 0) { + newGeom.move = glm::vec3(atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str())); + } utilityCore::safeGetline(fp_in, line); } diff --git a/src/sceneStructs.h b/src/sceneStructs.h index baa2e30..2856fd6 100644 --- a/src/sceneStructs.h +++ b/src/sceneStructs.h @@ -13,6 +13,12 @@ enum GeomType { struct Ray { glm::vec3 origin; glm::vec3 direction; + glm::vec3 color; + int index; + bool terminated; + int geomid; + bool out; + float time; }; struct Geom { @@ -24,6 +30,7 @@ struct Geom { glm::mat4 transform; glm::mat4 inverseTransform; glm::mat4 invTranspose; + glm::vec3 move; }; struct Material { diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu new file mode 100644 index 0000000..22f94a3 --- /dev/null +++ b/stream_compaction/efficient.cu @@ -0,0 +1,138 @@ +#include +#include +#include "common.h" +#include "efficient.h" + +namespace StreamCompaction { +namespace Efficient { + +int* g_odata; +int* g_idata; +int* dev_bools; +int* dev_indices; + +__global__ void generate_zeros(int *data) { + int i = threadIdx.x; + data[i] = 0; +} + +__global__ void set_zero(int size, int n, int *data) { + int i = threadIdx.x; + if (i >= n - 1) { + data[i] = 0; + } +} +__global__ void kern_up_sweep(int n, int *odata, const int *idata, int layer) { + int thrId = threadIdx.x + (blockIdx.x * blockDim.x); + if ((thrId < n) && (thrId%layer == 0)) { + odata[thrId + layer - 1] += idata[thrId + (layer / 2) - 1]; + } +} + +__global__ void kern_down_sweep(int n, int *odata, const int *idata, int layer) { + int thrId = threadIdx.x + (blockIdx.x * blockDim.x); + //__shared__ + if ((thrId < n) && (thrId%layer == 0)) { + int temp = idata[thrId + (layer / 2) - 1]; + odata[thrId + (layer / 2) - 1] = idata[thrId + layer - 1]; + odata[thrId + layer - 1] += temp; + } + + +} +/** + * Performs prefix-sum (aka scan) on idata, storing the result into odata. + */ +void scan(int n, int *odata, const int *idata) { + cudaEvent_t start, stop; + cudaEventCreate(&start); + cudaEventCreate(&stop); + + int blockSize = 128; + int numBlocks = ceil((float)n / (float)blockSize); + int powTwo = pow(2, ilog2ceil(n)); + dim3 fullBlocksPerGrid((powTwo + blockSize - 1) / blockSize); + cudaMalloc((void**)&g_odata, powTwo * sizeof(int)); + cudaMalloc((void**)&g_idata, powTwo * sizeof(int)); + + generate_zeros<<<1, powTwo>>>(g_odata); + generate_zeros<<<1, powTwo>>>(g_idata); + + int* scanArray = new int[n]; + //scanArray[0] = 0; + for (int i = 0; i < n; i++) { + scanArray[i] = idata[i]; + } + + cudaMemcpy(g_odata, odata, n*sizeof(int), cudaMemcpyHostToDevice); + cudaMemcpy(g_idata, scanArray, n*sizeof(int), cudaMemcpyHostToDevice); + + cudaEventRecord(start); + for (int d = 0; d <= ilog2ceil(n) - 1; d++) { + int layer = pow(2, d + 1); + g_odata = g_idata; + + kern_up_sweep<<>>(powTwo, g_odata, g_idata, layer); + g_idata = g_odata; + } + + + set_zero<<<1, powTwo>>>(powTwo, n, g_idata); + + for (int d = ilog2ceil(n) - 1; d >= 0; d--) { + int layer = pow(2, d + 1); + g_odata = g_idata; + kern_down_sweep<<>>(powTwo, g_odata, g_idata, layer); + g_idata = g_odata; + } + cudaEventRecord(stop); + + cudaEventSynchronize(stop); + float milliseconds = 0; + cudaEventElapsedTime(&milliseconds, start, stop); + //printf("%f - ", milliseconds); + cudaMemcpy(odata, g_odata, n*sizeof(int), cudaMemcpyDeviceToHost); +} + +/** + * Performs stream compaction on idata, storing the result into odata. + * All zeroes are discarded. + * + * @param n The number of elements in idata. + * @param odata The array into which to store elements. + * @param idata The array of elements to compact. + * @returns The number of elements remaining after compaction. + */ +int compact(int n, int *odata, const int *idata) { + int blockSize = 128; + int numBlocks = ceil((float)n / (float)blockSize); + int powTwo = pow(2, ilog2ceil(n)); + dim3 fullBlocksPerGrid((powTwo + blockSize - 1) / blockSize); + + cudaMalloc((void**)&g_odata, powTwo * sizeof(int)); + cudaMalloc((void**)&g_idata, powTwo * sizeof(int)); + cudaMalloc((void**)&dev_bools, powTwo * sizeof(int)); + cudaMalloc((void**)&dev_indices, powTwo * sizeof(int)); + + int* indices = new int[n]; + int* bools = new int[n]; + + cudaMemcpy(g_idata, idata, n*sizeof(int), cudaMemcpyHostToDevice); + + Common::kernMapToBoolean<<>>(powTwo, dev_bools, g_idata); + cudaMemcpy(bools, dev_bools, n*sizeof(int), cudaMemcpyDeviceToHost); + + scan(n, indices, bools); + + cudaMemcpy(dev_indices, indices, n*sizeof(int), cudaMemcpyHostToDevice); + + cudaMemcpy(g_idata, idata, n*sizeof(int), cudaMemcpyHostToDevice); + Common::kernScatter<<>>(powTwo, g_odata, g_idata, dev_bools, dev_indices); + + cudaMemcpy(odata, g_odata, n*sizeof(int), cudaMemcpyDeviceToHost); + + return indices[n-1] + bools[n-1]; +} + +} +} diff --git a/stream_compaction/efficient.h b/stream_compaction/efficient.h new file mode 100644 index 0000000..395ba10 --- /dev/null +++ b/stream_compaction/efficient.h @@ -0,0 +1,9 @@ +#pragma once + +namespace StreamCompaction { +namespace Efficient { + void scan(int n, int *odata, const int *idata); + + int compact(int n, int *odata, const int *idata); +} +}