diff --git a/CMakeLists.txt b/CMakeLists.txt index 76a65a0..4101fed 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -67,7 +67,7 @@ if(${CMAKE_SYSTEM_NAME} MATCHES "Darwin") set(CUDA_PROPAGATE_HOST_FLAGS OFF) endif() -#add_subdirectory(stream_compaction) # TODO: uncomment if using your stream compaction +add_subdirectory(stream_compaction) add_subdirectory(src) cuda_add_executable(${CMAKE_PROJECT_NAME} @@ -77,7 +77,7 @@ cuda_add_executable(${CMAKE_PROJECT_NAME} target_link_libraries(${CMAKE_PROJECT_NAME} src - #stream_compaction # TODO: uncomment if using your stream compaction + stream_compaction ${CORELIBS} ) diff --git a/Closed Scene Analysis.txt b/Closed Scene Analysis.txt new file mode 100644 index 0000000..1f84275 --- /dev/null +++ b/Closed Scene Analysis.txt @@ -0,0 +1,21 @@ +Closed Scene Analysis + +Stream compaction: +Depth 0: Trace took 4.87280ms and 476986 out of 640000 rays remain. +Depth 1: Trace took 3.80563ms and 403006 out of 640000 rays remain. +Depth 2: Trace took 3.32486ms and 348860 out of 640000 rays remain. +Depth 3: Trace took 2.91789ms and 305059 out of 640000 rays remain. +Depth 4: Trace took 2.58755ms and 268097 out of 640000 rays remain. +Depth 5: Trace took 2.29322ms and 236968 out of 640000 rays remain. +Depth 6: Trace took 2.05750ms and 209244 out of 640000 rays remain. +Depth 7: Trace took 1.83795ms and 185030 out of 640000 rays remain. + +No stream compaction: +Depth 0: Trace took 5.03702ms and 640000 out of 640000 rays remain. +Depth 1: Trace took 3.98282ms and 640000 out of 640000 rays remain. +Depth 2: Trace took 3.67200ms and 640000 out of 640000 rays remain. +Depth 3: Trace took 3.43274ms and 640000 out of 640000 rays remain. +Depth 4: Trace took 3.20074ms and 640000 out of 640000 rays remain. +Depth 5: Trace took 2.99098ms and 640000 out of 640000 rays remain. +Depth 6: Trace took 2.82339ms and 640000 out of 640000 rays remain. +Depth 7: Trace took 2.65965ms and 640000 out of 640000 rays remain. diff --git a/Open Scene Analysis.txt b/Open Scene Analysis.txt new file mode 100644 index 0000000..0ef560b --- /dev/null +++ b/Open Scene Analysis.txt @@ -0,0 +1,21 @@ +Open Scene Analysis + +Stream compaction: +Depth 0: Trace took 4.58256ms and 447407 out of 640000 rays remain. +Depth 1: Trace took 3.39939ms and 258788 out of 640000 rays remain. +Depth 2: Trace took 2.05389ms and 184010 out of 640000 rays remain. +Depth 3: Trace took 1.51395ms and 138638 out of 640000 rays remain. +Depth 4: Trace took 1.16778ms and 109147 out of 640000 rays remain. +Depth 5: Trace took 0.98333ms and 87256 out of 640000 rays remain. +Depth 6: Trace took 0.83344ms and 70458 out of 640000 rays remain. +Depth 7: Trace took 0.66774ms and 56732 out of 640000 rays remain. + +No stream compaction: +Depth 0: Trace took 4.66154ms and 640000 out of 640000 rays remain. +Depth 1: Trace took 3.53600ms and 640000 out of 640000 rays remain. +Depth 2: Trace took 2.62701ms and 640000 out of 640000 rays remain. +Depth 3: Trace took 2.29034ms and 640000 out of 640000 rays remain. +Depth 4: Trace took 2.06243ms and 640000 out of 640000 rays remain. +Depth 5: Trace took 1.90163ms and 640000 out of 640000 rays remain. +Depth 6: Trace took 1.77546ms and 640000 out of 640000 rays remain. +Depth 7: Trace took 1.66707ms and 640000 out of 640000 rays remain. diff --git a/README.md b/README.md index dbe1cbf..c3579cf 100644 --- a/README.md +++ b/README.md @@ -3,169 +3,90 @@ 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) +* Bradley Crusco +* Tested on: Windows 10, i7-3770K @ 3.50GHz 16GB, 2 x GTX 980 4096MB (Personal Computer) -### (TODO: Your README) +## Description +An interactive GPU accelerated path tracer with support for diffuse, specular, mirrored, and refractive surfaces. Additional features include depth of field and motion blur effects. +![](img/cornell_main_20k.png "Cornell Box") -*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. +Video of real time interaction: [https://youtu.be/Ja_5wvhphvI](https://youtu.be/Ja_5wvhphvI?rel=0) -Instructions (delete me) -======================== +## Features -This is due Thursday, September 24 evening at midnight. +### Diffuse Surfaces +![](img/cornell_diffuse.2015-09-27_02-29-25z.5000samp.png "Diffuse Sphere") -**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. +Diffuse surfaces are supported using a cosine weighted random direction calculation. -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.** +### Perfectly Specular Reflective Surfaces +![](img/cornell_mirror.2015-09-27_00-58-26z.5000samp.png "Perfectly Specular Mirror Sphere") -### Contents +Perfectly specular surfaces give a mirrored effect and are created by combining a specular light component with the calculation of the direction of a ray off a mirrored object. -* `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. +### Work Efficient Stream Compaction +Performance of the path tracer is optimized by using the work efficient stream compaction method found in [GPU Gems 3 Chapter 39: Parallel Prefix Sum (Scan) with CUDA](http://http.developer.nvidia.com/GPUGems3/gpugems3_ch39. html) on the array of rays. After each pass through all the rays, I check if a ray has been terminated (either because it hit a light or it traveled the maximum distance without intersecting with an object) and mark them as such. The stream compaction algorithm then takes this array and removes all of the terminated rays. This means at each pass through the rays, our collection of remaining ones to trace with becomes smaller, allowing us to free GPU threads for arrays that are still alive. More on this topic can be found in the [analysis](https://github.com/bcrusco/Project3-CUDA-Path-Tracer/blob/master/README.md#analysis) section bellow. +### Depth of Field +![](img/cornell_dof.2015-09-27_01-18-07z.5000_annotatedsamp.png "Depth of Field") -### Running the code +* **Overview**: When depth of field is disabled, my path tracer acts line a "pinhole camera". All the arrays come from a single point and are shot into each pixel of an image plane. Depth of field integrates over a lens to achieve its effect, dropping the pinhole implementation. I added two new configuration options to the camera in my scene files, focal distance and aperture radius. Focal distance specifies how far away from the camera is the image in focus, and replaces the idea of the image plane. And the aperture radius determines the severity of the effect (the blur of everything not at the focal distance). +* **Performance Impact**: Negligible. There is a few more calculations for depth of field than the standard pinhole implementation, but they are not major and they only take place when the rays are being created for the first bounce. So, the calculation will happen only once per iteration. +* **GPU vs. CPU Implementation**: The calculations for depth of field take place during the creation of my rays. This is the part of the code where the GPU and CPU implementations are most similar (the only difference being we can do it in parallel on the GPU). The specific depth of field calculations should be the same for either implementation, but it will just run faster on the GPU. -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`.) +### Non-Perfect Specular Surfaces +![](img/cornell_np_spec.2015-09-26_23-41-17z.5000samp.png "Non-Perfect Specular Surface") -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. +* **Overview**: Non-perfect specular surfaces, which gives a glossy effect, are created using a probability distribution between the diffuse and specular component of a material. First a probability of either a diffuse or a specular ray bounce occurring is calculating by weighting the intensity of the diffuse and specular color values respectively. A random value between 0 and 1 is then generated, which I use to choose a bounce type. The corresponding ray bounce direction is then calculated, as is the color, which is the given color provided by the scene file multiplied by the inverse probability that this bounce occurred. +* **Performance Impact**: Negligible. The only additional calculation to be done is the calculation of the ratio between both color intensities. There is a conditional, which may have performance impact, but this method only calculates one color and one ray bounce just like the mirrored and diffuse implementations. +* **GPU vs. CPU Implementation**: A CPU implementation would likely be recursive, where my GPU implementation is not. Because of this I use a probability calculation to determine how to bounce and only do the bounce once. Since the CPU implementation is recursive, it would likely trace both the specular and diffuse bounces instead of just picking one, and then use the ratio to determine the weights of the resulting color. So for the CPU implementation I would expect dramatically more performance requirements for this feature than my GPU implementation. -#### Controls +### Refractive Surfaces with Fresnel Effects +![](img/cornell_mirror.2015-09-27_00-58-26z.5000samp_annotated.png "Glass Sphere with Fresnel Effects") + +* **Overview**: This is calculated in much the same way as non-perfect specular surfaces. We figure out a probability that a ray hitting our refractive surface will either bounce off and reflect or pass into and refract through the object. If it reflects, we calculate the mirrored reflection direction, and if it refracts we calculate the ray direction using [Snell's law](https://en.wikipedia.org/wiki/Snell% 27s_law). The main difference is in the calculation of this probability. We calculate the Fresnel reflection coefficient using [Schlick's approximation](https://en.wikipedia.org/wiki/Schlick% 27s_approximation) (the inverse of which is the refraction coefficient). An index of refraction, specified in the scene files, determines the refractive properties of the respective material. Air has an index of refraction of 1, and glass about 2.2. It is important to keep track of whether a ray is going into an object or coming out of it, as the indexes are used in a ratio, and the ordering changes depending on what is being exited and what is being entered. +* **Performance Impact**: Significant. Compared to non-perfect specular surfaces, we have many more calculations to do to figure out the respective reflection and refraction coefficients. In addition, if a ray hits a refractive object at a perpendicular angle, the ray is always reflected, regardless of our reflection and refraction coefficients. This is another additional calculation that adds to the performance demands. +* **GPU vs. CPU Implementation**: As far as comparing my GPU implementation to what I'd expect a CPU implementation to be, it'd be the same as the comparison for non-perfect specular surfaces, except in this case the performance increase would be much more significant because it would have to make the additional calculations for the Fresnel coefficients. +* **How to Optimize**: To keep track of whether a ray was inside or outside of an object so I could know how to use the index of refraction coefficients (the air and the other object) I added a boolean to my Ray struct that held this state. This significantly added to my memory overhead because I'm using so many Ray's. I'd like to come up with a way to determine this property on the fly without saving it to further optimize performance. + +### Motion Blur +![](img/cornell_blur.2015-09-27_01-51-43z.5000samp.png "Motion Blur") +* **Overview**: Motion blur is very conceptually simple. We merely transform an object's position over the course of our render. This creates a blur effect because we are sampling the object at different locations, which creates a blurry trail as the object moves across the screen. Implementation was less trivial however, as originally I was not planning to support it and my geometry implementation did not support moving objects. To avoid changing a significant amount of code, I made a MovingGeom struct in addition to my original Geom struct to represent geometry that was moving. Since my path tracing implementation didn't have a concept of this MovingGeom, on every iteration before I begin path tracing I update the standard Geom of any object marked to be blurred, then trace as if it was static. The one additional change that was made to support this was a change to my scene files. It now requires "frame" tags to be added before transformation data. Two sets of data are expected, labeled "frame 0" and "frame 1", respectively. For objects that have motion blur enabled for them, these represent the starting and ending positions of the object. +* **Performance Impact**: Significant. The impact on the actual path tracing itself is nonexistent. There's no additional calculations, the objects we are intersecting with just happen to be in a different location. My specific workaround implementation to support a MovingGeom has significant consequences for memory bandwidth, as it demands that I load new Geom data on each iteration. Geoms can be very large, and are already a memory bottleneck, so this is less than ideal. If I want to eventually support arbitrary mesh models, I'll have to come up with a new implementation. +* **GPU vs. CPU Implementation**: The way I am calculating the motion blur effect is independent of how I am doing my path tracing, so there should be no difference between the CPU and GPU implementations. +* **How to Optimize**: As said above, the main space to optimize is the memory management of the MovingGeoms and Geoms. One more simple optimization, that is short of changing the entire project to support MovingGeoms, that might have a large effect is to store the MovingGeoms on the device memory instead of the host memory. As of now a transfer from host to device is required on every iteration, causing a big bottleneck. That bottleneck would at least be eliminated with this change. + + +## Analysis +### Stream Compaction: Open vs. Closed Scenes + +![](img/Project 3 Analysis 1.png "Active Threads Remaining at Trace Depth (Open vs. Closed Scene") +![](img/Project 3 Analysis 2.png "Trace Execution Time at Trace Depth (Open vs. Closed Scene") + +The two above charts compare an open cornell box (one of the walls is missing) vs. a closed box. The first thing we can see from these charts, especially the chart plotting active threads vs. trace depth, is that the stream compaction doesn't really kick in until the second bounce. For the open scene we see the biggest change between depth 0 to 1 and 1 to 2, before it starts to decrease less rapidly. It makes sense that we do not see too major of a drop off at index 0, because there hasn't yet been enough bounces for rays to reasonably terminate. + +Notice also how the closed scene data has a very gentle curve across the entire data set, compared to the sharp start off at the start for the open scene that then comes in line with the gentle curve. This shows a significant amount of rays that are bouncing out of the scene entirely. + +Regarding execution time, you can see that it correlates almost exactly with the active threads data. This shows us something fairly obvious, that it takes less time to trace over all the rays when there are less of them. This fact isn't very important here however, since the speedup is the result of losing rays that would otherwise contribute to our image. + +### Stream Compaction: Compaction vs. No Compaction + +![](img/Project 3 Analysis 3.png "Trace Execution Time at Trace Depth for an Open Scene (Compaction vs. No Compaction") +![](img/Project 3 Analysis 4.png "Trace Execution Time at Trace Depth for a Close Scene (Compaction vs. No Compaction") + +These above two charts compare the execution time of a trace across all remaining rays when stream compaction is enabled and disabled. In the case where stream compaction is disabled, in the ray bounce function itself I check, after it has already been launched, if it has been terminated. If so, then return. This stops the tracer from incorrectly calculating more bounces after the ray has been terminated, but keeps the overhead of the kernel launches that we are trying to avoid through compaction. + +The second chart, comparing compaction vs. no compaction for a closed scene, is the most important of the two, because the data in open scene chart is heavily influenced by the fact that rays are terminating early because they are leaving the scene through the open wall. Although the no compaction implementation isn't doing any more calculations than the stream compaction implementation, you can see that over time the gap between the two continues to widen. This is a great illustration of how the overhead of managing threads can drag down performance if it isn't managed correctly. Performance suffers because threads that actually need to execute will be waiting on launched threads that should be dead to figure out that they have nothing to do and free up GPU space for other rays. + +## Interactive 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. +* B activates and deactivates motion blur for a scene. Only works if at least one object in the scene is blur enabled. +* 0, activates and deactivates depth of field. 1 decreases focal distance, 2 increases. 3 decreases aperture radius, 3 increases. -## 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 -* 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) -* 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++. - -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); -``` - -### 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 +## 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 @@ -197,10 +118,13 @@ Cameras are defined in the following fashion: 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 -* frame (frame number) //start of a frame * 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 +* BLUR (bool blur) //motion blur flag for the entire scene, 0 for no, 1 for yes +* DOF (bool dof) //depth of field flag, 0 for no, 1 for yes +* FD (float fd) //focal distance for depth of field +* AR (float ar) //aperture radius for depth of field Objects are defined in the following fashion: @@ -209,75 +133,8 @@ Objects are defined in the following fashion: "mesh". Note that cubes and spheres are unit sized and centered at the origin. * material (material ID) //material to assign this object -* frame (frame number) //start of a frame +* BLUR (bool blur) //motion blur flag for individual object, 0 for no, 1 for yes +* frame (int framenum) //0 or 1. Used for motion blur. Specify start and end transforms * 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! - -### Analysis - -* Stream compaction helps most after a few bounces. Print and plot the - 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. -* 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? - - -## 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/AA comparison.psd b/img/AA comparison.psd new file mode 100644 index 0000000..c99a3ff Binary files /dev/null and b/img/AA comparison.psd differ diff --git a/img/Project 3 Analysis 1.png b/img/Project 3 Analysis 1.png new file mode 100644 index 0000000..6ede7e5 Binary files /dev/null and b/img/Project 3 Analysis 1.png differ diff --git a/img/Project 3 Analysis 2.png b/img/Project 3 Analysis 2.png new file mode 100644 index 0000000..94dfb9f Binary files /dev/null and b/img/Project 3 Analysis 2.png differ diff --git a/img/Project 3 Analysis 3.png b/img/Project 3 Analysis 3.png new file mode 100644 index 0000000..67b2cb6 Binary files /dev/null and b/img/Project 3 Analysis 3.png differ diff --git a/img/Project 3 Analysis 4.png b/img/Project 3 Analysis 4.png new file mode 100644 index 0000000..8a6c902 Binary files /dev/null and b/img/Project 3 Analysis 4.png differ diff --git a/img/REFERENCE_cornell.5000samp.png b/img/REFERENCE_cornell.5000samp.png deleted file mode 100644 index 5ceb26e..0000000 Binary files a/img/REFERENCE_cornell.5000samp.png and /dev/null differ diff --git a/img/REFERENCE_sphere.5000samp.png b/img/REFERENCE_sphere.5000samp.png deleted file mode 100644 index 0628927..0000000 Binary files a/img/REFERENCE_sphere.5000samp.png and /dev/null differ diff --git a/img/cornell_blur.2015-09-27_01-51-43z.5000samp.png b/img/cornell_blur.2015-09-27_01-51-43z.5000samp.png new file mode 100644 index 0000000..8e84e09 Binary files /dev/null and b/img/cornell_blur.2015-09-27_01-51-43z.5000samp.png differ diff --git a/img/cornell_blur.2015-10-02_15-48-32z.20000samp.png b/img/cornell_blur.2015-10-02_15-48-32z.20000samp.png new file mode 100644 index 0000000..07bcf0a Binary files /dev/null and b/img/cornell_blur.2015-10-02_15-48-32z.20000samp.png differ diff --git a/img/cornell_diffuse.2015-09-27_02-29-25z.5000samp.png b/img/cornell_diffuse.2015-09-27_02-29-25z.5000samp.png new file mode 100644 index 0000000..e337d51 Binary files /dev/null and b/img/cornell_diffuse.2015-09-27_02-29-25z.5000samp.png differ diff --git a/img/cornell_dof.2015-09-27_01-18-07z.5000_annotatedsamp.png b/img/cornell_dof.2015-09-27_01-18-07z.5000_annotatedsamp.png new file mode 100644 index 0000000..a6801bd Binary files /dev/null and b/img/cornell_dof.2015-09-27_01-18-07z.5000_annotatedsamp.png differ diff --git a/img/cornell_dof.2015-09-27_01-18-07z.5000samp.png b/img/cornell_dof.2015-09-27_01-18-07z.5000samp.png new file mode 100644 index 0000000..87cd2ed Binary files /dev/null and b/img/cornell_dof.2015-09-27_01-18-07z.5000samp.png differ diff --git a/img/cornell_glass.2015-10-02_15-13-50z.20000samp.png b/img/cornell_glass.2015-10-02_15-13-50z.20000samp.png new file mode 100644 index 0000000..a7d1f8e Binary files /dev/null and b/img/cornell_glass.2015-10-02_15-13-50z.20000samp.png differ diff --git a/img/cornell_glass_10k.png b/img/cornell_glass_10k.png new file mode 100644 index 0000000..a64e3f5 Binary files /dev/null and b/img/cornell_glass_10k.png differ diff --git a/img/cornell_main_20k.png b/img/cornell_main_20k.png new file mode 100644 index 0000000..00ac679 Binary files /dev/null and b/img/cornell_main_20k.png differ diff --git a/img/cornell_mirror.2015-09-27_00-58-26z.5000samp.png b/img/cornell_mirror.2015-09-27_00-58-26z.5000samp.png new file mode 100644 index 0000000..e22d160 Binary files /dev/null and b/img/cornell_mirror.2015-09-27_00-58-26z.5000samp.png differ diff --git a/img/cornell_mirror.2015-09-27_00-58-26z.5000samp_annotated.png b/img/cornell_mirror.2015-09-27_00-58-26z.5000samp_annotated.png new file mode 100644 index 0000000..8d56502 Binary files /dev/null and b/img/cornell_mirror.2015-09-27_00-58-26z.5000samp_annotated.png differ diff --git a/img/cornell_noaa.2015-09-27_02-02-48z.5000samp.png b/img/cornell_noaa.2015-09-27_02-02-48z.5000samp.png new file mode 100644 index 0000000..311917f Binary files /dev/null and b/img/cornell_noaa.2015-09-27_02-02-48z.5000samp.png differ diff --git a/img/cornell_np_spec.2015-09-26_23-41-17z.5000samp.png b/img/cornell_np_spec.2015-09-26_23-41-17z.5000samp.png new file mode 100644 index 0000000..90d5a9b Binary files /dev/null and b/img/cornell_np_spec.2015-09-26_23-41-17z.5000samp.png differ diff --git a/scenes/cornell.txt b/scenes/cornell.txt index 7c71b63..5dc1716 100644 --- a/scenes/cornell.txt +++ b/scenes/cornell.txt @@ -6,10 +6,7 @@ SPECRGB 0 0 0 REFL 0 REFR 0 REFRIOR 0 -SCATTER 0 -ABSCOEFF 0 0 0 -RSCTCOEFF 0 -EMITTANCE 5 +EMITTANCE 1 // Diffuse white MATERIAL 1 @@ -19,9 +16,6 @@ SPECRGB 1 1 1 REFL 0 REFR 0 REFRIOR 0 -SCATTER 0 -ABSCOEFF 0 0 0 -RSCTCOEFF 0 EMITTANCE 0 // Diffuse red @@ -32,9 +26,6 @@ SPECRGB 1 1 1 REFL 0 REFR 0 REFRIOR 0 -SCATTER 0 -ABSCOEFF 0 0 0 -RSCTCOEFF 0 EMITTANCE 0 // Diffuse green @@ -45,83 +36,251 @@ SPECRGB 1 1 1 REFL 0 REFR 0 REFRIOR 0 -SCATTER 0 -ABSCOEFF 0 0 0 -RSCTCOEFF 0 +EMITTANCE 0 + +// spec green +MATERIAL 4 +RGB 0 1 0 +SPECEX .01 +SPECRGB .2 .2 .2 +REFL 1 +REFR 0 +REFRIOR 1.5 +EMITTANCE 0 + +// glass +MATERIAL 5 +RGB .25 .85 .35 +SPECEX 1 +SPECRGB 1 1 1 +REFL 0 +REFR 1 +REFRIOR 2.2 +EMITTANCE 0 + +// spec blue +MATERIAL 6 +RGB .11 .56 1 +SPECEX .01 +SPECRGB .1 .1 .1 +REFL 1 +REFR 0 +REFRIOR 1.5 +EMITTANCE 0 + +// Perfect mirror +MATERIAL 7 +RGB 1 1 1 +SPECEX 0 +SPECRGB 1 1 1 +REFL 1 +REFR 0 +REFRIOR 1.5 +EMITTANCE 0 + +// spec red +MATERIAL 8 +RGB 1 .18 .18 +SPECEX .01 +SPECRGB .1 .1 .1 +REFL 1 +REFR 0 +REFRIOR 1.5 EMITTANCE 0 // Camera CAMERA RES 800 800 FOVY 45 -ITERATIONS 5000 +ITERATIONS 20000 DEPTH 8 -FILE cornell -frame 0 +FILE cornell_mirror EYE 0.0 5 10.5 VIEW 0 0 -1 UP 0 1 0 - +BLUR 0 // Is blur active +DOF 0 // Is DOF active +FD 9.5 // Focal distance +AR 0.2 // Aperture radius // Ceiling light OBJECT 0 cube material 0 +BLUR 0 frame 0 TRANS 0 10 0 ROTAT 0 0 0 -SCALE 3 .3 3 +SCALE 10 .2 30 +frame 1 +TRANS 0 10 0 +ROTAT 0 0 0 +SCALE 10 .2 30 // Floor OBJECT 1 cube material 1 +BLUR 0 frame 0 TRANS 0 0 0 ROTAT 0 0 0 -SCALE 10 .01 10 +SCALE 10 .01 30 +frame 1 +TRANS 0 0 0 +ROTAT 0 0 0 +SCALE 10 .01 30 // Ceiling OBJECT 2 cube material 1 +BLUR 0 frame 0 TRANS 0 10 0 ROTAT 0 0 90 -SCALE .01 10 10 +SCALE .01 10 30 +frame 1 +TRANS 0 10 0 +ROTAT 0 0 90 +SCALE .01 10 30 // Back wall OBJECT 3 cube material 1 +BLUR 0 frame 0 TRANS 0 5 -5 ROTAT 0 90 0 SCALE .01 10 10 +frame 1 +TRANS 0 5 -5 +ROTAT 0 90 0 +SCALE .01 10 10 // Left wall OBJECT 4 cube material 2 +BLUR 0 frame 0 TRANS -5 5 0 ROTAT 0 0 0 -SCALE .01 10 10 +SCALE .01 10 30 +frame 1 +TRANS -5 5 0 +ROTAT 0 0 0 +SCALE .01 10 30 // Right wall OBJECT 5 cube material 3 +BLUR 0 frame 0 TRANS 5 5 0 ROTAT 0 0 0 -SCALE .01 10 10 +SCALE .01 10 30 +frame 1 +TRANS 5 5 0 +ROTAT 0 0 0 +SCALE .01 10 30 // Sphere OBJECT 6 sphere -material 1 +material 4 +BLUR 0 +frame 0 +TRANS -2 5 -1 +ROTAT 0 0 0 +SCALE 3 3 3 +frame 1 +TRANS -2 5 -1 +ROTAT 0 0 0 +SCALE 3 3 3 + +// Glass box +OBJECT 7 +cube +material 5 +BLUR 0 +frame 0 +TRANS 2 1 -1 +ROTAT 0 45 0 +SCALE 3 3 3 +frame 1 +TRANS 2 1 -1 +ROTAT 0 45 0 +SCALE 3 3 3 + +// sphere +OBJECT 8 +sphere +material 6 +BLUR 0 +frame 0 +TRANS 2 3.5 -1 +ROTAT 0 0 0 +SCALE 2 2 2 +frame 1 +TRANS 2 3.5 -1 +ROTAT 0 0 0 +SCALE 2 2 2 + +// Perfect Mirror box +OBJECT 9 +cube +material 7 +BLUR 0 frame 0 -TRANS -1 4 -1 +TRANS -3 1 0 +ROTAT 0 60 0 +SCALE 2 2 2 +frame 1 +TRANS -3 1 0 +ROTAT 0 60 0 +SCALE 2 2 2 + +// non perfect red sphere +OBJECT 10 +sphere +material 8 +BLUR 0 +frame 0 +TRANS 2.5 6.5 -2.8 +ROTAT 0 0 0 +SCALE 3 3 3 +frame 1 +TRANS 2.5 6.5 -2.8 ROTAT 0 0 0 SCALE 3 3 3 + +// glass ball +OBJECT 11 +sphere +material 5 +BLUR 0 +frame 0 +TRANS 3.8 .7 .9 +ROTAT 0 0 0 +SCALE 1.5 1.5 1.5 +frame 1 +TRANS 3.8 .7 .9 +ROTAT 0 0 0 +SCALE 1.5 1.5 1.5 + +// front wall +OBJECT 12 +cube +material 1 +BLUR 0 +frame 0 +TRANS 0 5 15 +ROTAT 0 90 0 +SCALE .01 10 10 +frame 1 +TRANS 0 5 15 +ROTAT 0 90 0 +SCALE .01 10 10 diff --git a/scenes/cornell_blur_simple.txt b/scenes/cornell_blur_simple.txt new file mode 100644 index 0000000..bd1920f --- /dev/null +++ b/scenes/cornell_blur_simple.txt @@ -0,0 +1,162 @@ +// Emissive material (light) +MATERIAL 0 +RGB 1 1 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 1 + +// 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 + +// spec green +MATERIAL 4 +RGB 0 1 0 +SPECEX .01 +SPECRGB .2 .2 .2 +REFL 1 +REFR 0 +REFRIOR 1.5 +EMITTANCE 0 + +// Camera +CAMERA +RES 800 800 +FOVY 45 +ITERATIONS 5000 +DEPTH 8 +FILE cornell_blur +EYE 0.0 5 10.5 +VIEW 0 0 -1 +UP 0 1 0 +BLUR 1 // Is blur active +DOF 0 // Is DOF active +FD 9.5 // Focal distance +AR 0.2 // Aperture radius + +// Ceiling light +OBJECT 0 +cube +material 0 +BLUR 0 +frame 0 +TRANS 0 10 0 +ROTAT 0 0 0 +SCALE 10 .2 10 +frame 1 +TRANS 0 10 0 +ROTAT 0 0 0 +SCALE 10 .2 10 + +// Floor +OBJECT 1 +cube +material 1 +BLUR 0 +frame 0 +TRANS 0 0 0 +ROTAT 0 0 0 +SCALE 10 .01 10 +frame 1 +TRANS 0 0 0 +ROTAT 0 0 0 +SCALE 10 .01 10 + +// Ceiling +OBJECT 2 +cube +material 1 +BLUR 0 +frame 0 +TRANS 0 10 0 +ROTAT 0 0 90 +SCALE .01 10 10 +frame 1 +TRANS 0 10 0 +ROTAT 0 0 90 +SCALE .01 10 10 + +// Back wall +OBJECT 3 +cube +material 1 +BLUR 0 +frame 0 +TRANS 0 5 -5 +ROTAT 0 90 0 +SCALE .01 10 10 +frame 1 +TRANS 0 5 -5 +ROTAT 0 90 0 +SCALE .01 10 10 + +// Left wall +OBJECT 4 +cube +material 2 +BLUR 0 +frame 0 +TRANS -5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 +frame 1 +TRANS -5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +// Right wall +OBJECT 5 +cube +material 3 +BLUR 0 +frame 0 +TRANS 5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 +frame 1 +TRANS 5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +// Sphere +OBJECT 6 +sphere +material 4 +BLUR 1 +frame 0 +TRANS -2 4 -1 +ROTAT 0 0 0 +SCALE 3 3 3 +frame 1 +TRANS -4 4 -1 +ROTAT 0 0 0 +SCALE 3 3 3 diff --git a/scenes/cornell_closed.txt b/scenes/cornell_closed.txt new file mode 100644 index 0000000..23b45b3 --- /dev/null +++ b/scenes/cornell_closed.txt @@ -0,0 +1,288 @@ +// Emissive material (light) +MATERIAL 0 +RGB 1 1 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 1 + +// 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 + +// spec green +MATERIAL 4 +RGB 0 1 0 +SPECEX .01 +SPECRGB .2 .2 .2 +REFL 1 +REFR 0 +REFRIOR 1.5 +EMITTANCE 0 + +// glass +MATERIAL 5 +RGB .25 .85 .35 +SPECEX 1 +SPECRGB 1 1 1 +REFL 0 +REFR 1 +REFRIOR 2.2 +EMITTANCE 0 + +// spec blue +MATERIAL 6 +RGB .11 .56 1 +SPECEX .01 +SPECRGB .1 .1 .1 +REFL 1 +REFR 0 +REFRIOR 1.5 +EMITTANCE 0 + +// Perfect mirror +MATERIAL 7 +RGB 1 1 1 +SPECEX 0 +SPECRGB 1 1 1 +REFL 1 +REFR 0 +REFRIOR 1.5 +EMITTANCE 0 + +// diffuse gold +MATERIAL 8 +RGB 1 .84 0 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 1.5 +EMITTANCE 0 + +// Camera +CAMERA +RES 800 800 +FOVY 45 +ITERATIONS 5000 +DEPTH 8 +FILE cornell_mirror +EYE 0.0 5 4.5 +VIEW 0 0 -1 +UP 0 1 0 +BLUR 0 // Is blur active +DOF 0 // Is DOF active +FD 9.5 // Focal distance +AR 0.2 // Aperture radius + +// Ceiling light +OBJECT 0 +cube +material 0 +BLUR 0 +frame 0 +TRANS 0 10 0 +ROTAT 0 0 0 +SCALE 10 3 10 +frame 1 +TRANS 0 10 0 +ROTAT 0 0 0 +SCALE 10 3 10 + +// Floor +OBJECT 1 +cube +material 1 +BLUR 0 +frame 0 +TRANS 0 0 0 +ROTAT 0 0 0 +SCALE 10 .01 10 +frame 1 +TRANS 0 0 0 +ROTAT 0 0 0 +SCALE 10 .01 10 + +// Ceiling +OBJECT 2 +cube +material 1 +BLUR 0 +frame 0 +TRANS 0 10 0 +ROTAT 0 0 90 +SCALE .01 10 10 +frame 1 +TRANS 0 10 0 +ROTAT 0 0 90 +SCALE .01 10 10 + + +// Back wall +OBJECT 3 +cube +material 1 +BLUR 0 +frame 0 +TRANS 0 5 -5 +ROTAT 0 90 0 +SCALE .01 10 10 +frame 1 +TRANS 0 5 -5 +ROTAT 0 90 0 +SCALE .01 10 10 + +// Left wall +OBJECT 4 +cube +material 2 +BLUR 0 +frame 0 +TRANS -5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 +frame 1 +TRANS -5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +// Right wall +OBJECT 5 +cube +material 3 +BLUR 0 +frame 0 +TRANS 5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 +frame 1 +TRANS 5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +// Sphere +OBJECT 6 +sphere +material 4 +BLUR 0 +frame 0 +TRANS -2 5 -1 +ROTAT 0 0 0 +SCALE 3 3 3 +frame 1 +TRANS -2 5 -1 +ROTAT 0 0 0 +SCALE 3 3 3 + +// Glass box +OBJECT 7 +cube +material 5 +BLUR 0 +frame 0 +TRANS 2 1 -1 +ROTAT 0 45 0 +SCALE 3 3 3 +frame 1 +TRANS 2 1 -1 +ROTAT 0 45 0 +SCALE 3 3 3 + +// sphere +OBJECT 8 +sphere +material 6 +BLUR 0 +frame 0 +TRANS 2 3.5 -1 +ROTAT 0 0 0 +SCALE 2 2 2 +frame 1 +TRANS 2 3.5 -1 +ROTAT 0 0 0 +SCALE 2 2 2 + +// Perfect Mirror box +OBJECT 9 +cube +material 7 +BLUR 0 +frame 0 +TRANS -3 1 0 +ROTAT 0 60 0 +SCALE 2 2 2 +frame 1 +TRANS -3 1 0 +ROTAT 0 60 0 +SCALE 2 2 2 + +// diffuse sphere +OBJECT 10 +sphere +material 8 +BLUR 0 +frame 0 +TRANS 2.5 6.5 -2.8 +ROTAT 0 0 0 +SCALE 3 3 3 +frame 1 +TRANS 2.5 6.5 -2.8 +ROTAT 0 0 0 +SCALE 3 3 3 + +// glass ball +OBJECT 11 +sphere +material 5 +BLUR 0 +frame 0 +TRANS 3.8 .7 .9 +ROTAT 0 0 0 +SCALE 1.5 1.5 1.5 +frame 1 +TRANS 3.8 .7 .9 +ROTAT 0 0 0 +SCALE 1.5 1.5 1.5 + + +// front wall +OBJECT 12 +cube +material 1 +BLUR 0 +frame 0 +TRANS 0 5 5 +ROTAT 0 90 0 +SCALE .01 10 10 +frame 1 +TRANS 0 5 5 +ROTAT 0 90 0 +SCALE .01 10 10 diff --git a/scenes/cornell_glass_simple.txt b/scenes/cornell_glass_simple.txt new file mode 100644 index 0000000..94aa069 --- /dev/null +++ b/scenes/cornell_glass_simple.txt @@ -0,0 +1,162 @@ +// Emissive material (light) +MATERIAL 0 +RGB 1 1 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 10 + +// 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 + +// spec green +MATERIAL 4 +RGB .25 .85 .35 +SPECEX 1 +SPECRGB 1 1 1 +REFL 0 +REFR 1 +REFRIOR 2.2 +EMITTANCE 0 + +// Camera +CAMERA +RES 800 800 +FOVY 45 +ITERATIONS 5000 +DEPTH 8 +FILE cornell_glass +EYE 0.0 5 10.5 +VIEW 0 0 -1 +UP 0 1 0 +BLUR 1 // Is blur active +DOF 0 // Is DOF active +FD 9.5 // Focal distance +AR 0.2 // Aperture radius + +// Ceiling light +OBJECT 0 +cube +material 0 +BLUR 0 +frame 0 +TRANS 0 10 0 +ROTAT 0 0 0 +SCALE 3 0.2 3 +frame 1 +TRANS 0 10 0 +ROTAT 0 0 0 +SCALE 3 0.2 3 + +// Floor +OBJECT 1 +cube +material 1 +BLUR 0 +frame 0 +TRANS 0 0 0 +ROTAT 0 0 0 +SCALE 10 .01 10 +frame 1 +TRANS 0 0 0 +ROTAT 0 0 0 +SCALE 10 .01 10 + +// Ceiling +OBJECT 2 +cube +material 1 +BLUR 0 +frame 0 +TRANS 0 10 0 +ROTAT 0 0 90 +SCALE .01 10 10 +frame 1 +TRANS 0 10 0 +ROTAT 0 0 90 +SCALE .01 10 10 + +// Back wall +OBJECT 3 +cube +material 1 +BLUR 0 +frame 0 +TRANS 0 5 -5 +ROTAT 0 90 0 +SCALE .01 10 10 +frame 1 +TRANS 0 5 -5 +ROTAT 0 90 0 +SCALE .01 10 10 + +// Left wall +OBJECT 4 +cube +material 2 +BLUR 0 +frame 0 +TRANS -5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 +frame 1 +TRANS -5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +// Right wall +OBJECT 5 +cube +material 3 +BLUR 0 +frame 0 +TRANS 5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 +frame 1 +TRANS 5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +// Sphere +OBJECT 6 +sphere +material 4 +BLUR 0 +frame 0 +TRANS -1 4 -1 +ROTAT 0 0 0 +SCALE 3 3 3 +frame 1 +TRANS -4 4 -1 +ROTAT 0 0 0 +SCALE 3 3 3 diff --git a/scenes/cornell_mirror_simple.txt b/scenes/cornell_mirror_simple.txt new file mode 100644 index 0000000..412973a --- /dev/null +++ b/scenes/cornell_mirror_simple.txt @@ -0,0 +1,162 @@ +// Emissive material (light) +MATERIAL 0 +RGB 1 1 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 1 + +// 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 + +// spec blue +MATERIAL 4 +RGB .11 .56 1 +SPECEX .01 +SPECRGB .1 .1 .1 +REFL 1 +REFR 0 +REFRIOR 1.5 +EMITTANCE 0 + +// Camera +CAMERA +RES 800 800 +FOVY 45 +ITERATIONS 5000 +DEPTH 8 +FILE cornell_mirror +EYE 0.0 5 10.5 +VIEW 0 0 -1 +UP 0 1 0 +BLUR 0 // Is blur active +DOF 0 // Is DOF active +FD 9.5 // Focal distance +AR 0.2 // Aperture radius + +// Ceiling light +OBJECT 0 +cube +material 0 +BLUR 0 +frame 0 +TRANS 0 10 0 +ROTAT 0 0 0 +SCALE 10 .2 10 +frame 1 +TRANS 0 10 0 +ROTAT 0 0 0 +SCALE 10 .2 10 + +// Floor +OBJECT 1 +cube +material 1 +BLUR 0 +frame 0 +TRANS 0 0 0 +ROTAT 0 0 0 +SCALE 10 .01 10 +frame 1 +TRANS 0 0 0 +ROTAT 0 0 0 +SCALE 10 .01 10 + +// Ceiling +OBJECT 2 +cube +material 1 +BLUR 0 +frame 0 +TRANS 0 10 0 +ROTAT 0 0 90 +SCALE .01 10 10 +frame 1 +TRANS 0 10 0 +ROTAT 0 0 90 +SCALE .01 10 10 + +// Back wall +OBJECT 3 +cube +material 1 +BLUR 0 +frame 0 +TRANS 0 5 -5 +ROTAT 0 90 0 +SCALE .01 10 10 +frame 1 +TRANS 0 5 -5 +ROTAT 0 90 0 +SCALE .01 10 10 + +// Left wall +OBJECT 4 +cube +material 2 +BLUR 0 +frame 0 +TRANS -5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 +frame 1 +TRANS -5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +// Right wall +OBJECT 5 +cube +material 3 +BLUR 0 +frame 0 +TRANS 5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 +frame 1 +TRANS 5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +// Sphere +OBJECT 6 +sphere +material 4 +BLUR 0 +frame 0 +TRANS -1 4 -1 +ROTAT 0 0 0 +SCALE 3 3 3 +frame 1 +TRANS -1 4 -1 +ROTAT 0 0 0 +SCALE 3 3 3 diff --git a/scenes/cornell_perfect_mirror_simple.txt b/scenes/cornell_perfect_mirror_simple.txt new file mode 100644 index 0000000..9581635 --- /dev/null +++ b/scenes/cornell_perfect_mirror_simple.txt @@ -0,0 +1,162 @@ +// Emissive material (light) +MATERIAL 0 +RGB 1 1 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 1 + +// 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 1 +REFR 0 +REFRIOR 1.5 +EMITTANCE 0 + +// Camera +CAMERA +RES 800 800 +FOVY 45 +ITERATIONS 5000 +DEPTH 8 +FILE cornell_mirror +EYE 0.0 5 10.5 +VIEW 0 0 -1 +UP 0 1 0 +BLUR 0 // Is blur active +DOF 0 // Is DOF active +FD 9.5 // Focal distance +AR 0.2 // Aperture radius + +// Ceiling light +OBJECT 0 +cube +material 0 +BLUR 0 +frame 0 +TRANS 0 10 0 +ROTAT 0 0 0 +SCALE 10 .2 10 +frame 1 +TRANS 0 10 0 +ROTAT 0 0 0 +SCALE 10 .2 10 + +// Floor +OBJECT 1 +cube +material 1 +BLUR 0 +frame 0 +TRANS 0 0 0 +ROTAT 0 0 0 +SCALE 10 .01 10 +frame 1 +TRANS 0 0 0 +ROTAT 0 0 0 +SCALE 10 .01 10 + +// Ceiling +OBJECT 2 +cube +material 1 +BLUR 0 +frame 0 +TRANS 0 10 0 +ROTAT 0 0 90 +SCALE .01 10 10 +frame 1 +TRANS 0 10 0 +ROTAT 0 0 90 +SCALE .01 10 10 + +// Back wall +OBJECT 3 +cube +material 1 +BLUR 0 +frame 0 +TRANS 0 5 -5 +ROTAT 0 90 0 +SCALE .01 10 10 +frame 1 +TRANS 0 5 -5 +ROTAT 0 90 0 +SCALE .01 10 10 + +// Left wall +OBJECT 4 +cube +material 2 +BLUR 0 +frame 0 +TRANS -5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 +frame 1 +TRANS -5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +// Right wall +OBJECT 5 +cube +material 3 +BLUR 0 +frame 0 +TRANS 5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 +frame 1 +TRANS 5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +// Sphere +OBJECT 6 +sphere +material 4 +BLUR 0 +frame 0 +TRANS -1 4 -1 +ROTAT 0 0 0 +SCALE 3 3 3 +frame 1 +TRANS -1 4 -1 +ROTAT 0 0 0 +SCALE 3 3 3 diff --git a/scenes/cornell_simple.txt b/scenes/cornell_simple.txt new file mode 100644 index 0000000..79b87eb --- /dev/null +++ b/scenes/cornell_simple.txt @@ -0,0 +1,162 @@ +// Emissive material (light) +MATERIAL 0 +RGB 1 1 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 1 + +// 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 + +// spec green +MATERIAL 4 +RGB 0 1 0 +SPECEX .01 +SPECRGB .2 .2 .2 +REFL 1 +REFR 0 +REFRIOR 1.5 +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 +BLUR 0 // Is blur active +DOF 0 // Is DOF active +FD 9.5 // Focal distance +AR 0.2 // Aperture radius + +// Ceiling light +OBJECT 0 +cube +material 0 +BLUR 0 +frame 0 +TRANS 0 10 0 +ROTAT 0 0 0 +SCALE 10 3 10 +frame 1 +TRANS 0 10 0 +ROTAT 0 0 0 +SCALE 10 3 10 + +// Floor +OBJECT 1 +cube +material 1 +BLUR 0 +frame 0 +TRANS 0 0 0 +ROTAT 0 0 0 +SCALE 10 .01 10 +frame 1 +TRANS 0 0 0 +ROTAT 0 0 0 +SCALE 10 .01 10 + +// Ceiling +OBJECT 2 +cube +material 1 +BLUR 0 +frame 0 +TRANS 0 10 0 +ROTAT 0 0 90 +SCALE .01 10 10 +frame 1 +TRANS 0 10 0 +ROTAT 0 0 90 +SCALE .01 10 10 + +// Back wall +OBJECT 3 +cube +material 1 +BLUR 0 +frame 0 +TRANS 0 5 -5 +ROTAT 0 90 0 +SCALE .01 10 10 +frame 1 +TRANS 0 5 -5 +ROTAT 0 90 0 +SCALE .01 10 10 + +// Left wall +OBJECT 4 +cube +material 2 +BLUR 0 +frame 0 +TRANS -5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 +frame 1 +TRANS -5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +// Right wall +OBJECT 5 +cube +material 3 +BLUR 0 +frame 0 +TRANS 5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 +frame 1 +TRANS 5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +// Sphere +OBJECT 6 +sphere +material 1 +BLUR 0 +frame 0 +TRANS -1 4 -1 +ROTAT 0 0 0 +SCALE 3 3 3 +frame 1 +TRANS -1 4 -1 +ROTAT 0 0 0 +SCALE 3 3 3 diff --git a/scenes/sphere.txt b/scenes/sphere.txt index 20149a2..c70d546 100644 --- a/scenes/sphere.txt +++ b/scenes/sphere.txt @@ -6,9 +6,6 @@ SPECRGB 0 0 0 REFL 0 REFR 0 REFRIOR 0 -SCATTER 0 -ABSCOEFF 0 0 0 -RSCTCOEFF 0 EMITTANCE 5 // Camera @@ -18,7 +15,6 @@ FOVY 45 ITERATIONS 5000 DEPTH 8 FILE sphere -frame 0 EYE 0.0 5 10.5 VIEW 0 0 -1 UP 0 1 0 @@ -27,7 +23,6 @@ UP 0 1 0 OBJECT 0 sphere material 0 -frame 0 TRANS 0 0 0 ROTAT 0 0 0 SCALE 3 3 3 diff --git a/src/interactions.h b/src/interactions.h index 22c1710..fef130b 100644 --- a/src/interactions.h +++ b/src/interactions.h @@ -31,38 +31,159 @@ glm::vec3 calculateRandomDirectionInHemisphere( } // Use not-normal direction to generate two perpendicular directions - glm::vec3 perpendicularDirection1 = + glm::vec3 perpendicularDirection = glm::normalize(glm::cross(normal, directionNotNormal)); - glm::vec3 perpendicularDirection2 = - glm::normalize(glm::cross(normal, perpendicularDirection1)); return up * normal - + cos(around) * over * perpendicularDirection1 - + sin(around) * over * perpendicularDirection2; + + cos(around) * over * perpendicularDirection + + sin(around) * over * (glm::normalize(glm::cross(normal, perpendicularDirection))); } /** - * Scatter a ray with some probabilities according to the material properties. - * For example, a diffuse surface scatters in a cosine-weighted hemisphere. - * A perfect specular surface scatters in the reflected ray direction. - * In order to apply multiple effects to one surface, probabilistically choose - * between them. - * - * This method applies its changes to the Ray parameter `ray` in place. - * It also modifies the color `color` of the ray in place. - * - * You may need to change the parameter list for your purposes! - */ +* Computes Fresnel reflection coefficient using Schlick's approximation +* https://en.wikipedia.org/wiki/Schlick%27s_approximation +* http://www.scratchapixel.com/old/lessons/3d-basic-lessons/lesson-14-interaction-light-matter/optics-reflection-and-refraction/ +*/ +__host__ __device__ +float calculateFresnelReflectionCoefficient(glm::vec3 direction, glm::vec3 normal, float indexOfRefraction) { + float r0 = glm::pow((1.0f - indexOfRefraction) / (1.0f + indexOfRefraction), 2); + return r0 + (1.0f - r0) * glm::pow(1.0f - glm::dot(normal, -direction), 5); +} + +/** +* Computes a reflection vector off a specular or refractive surface. +* Used for specular and refracted lighting. +*/ +__host__ __device__ +glm::vec3 calculateReflectionDirection(glm::vec3 direction, glm::vec3 normal) { + return direction + 2.0f * glm::dot(-direction, normal) * normal; +} + +/** +* Computes a refraction vector off a refractive surface. +* Used for refracted lighting. +*/ +__host__ __device__ +glm::vec3 calculateRefractionDirection(glm::vec3 direction, glm::vec3 normal, float angle, float eta) { + return (-eta * glm::dot(normal, direction) - glm::sqrt(angle)) * normal + direction * eta; +} + +/** +* Updates the transformation of a geom to move over the duration of the render from configured start and end locations. +* Right now only supports translations. +*/ +__host__ __device__ +void motionBlur(MovingGeom *mgeoms, int id, int iter, int maxIter) { + if (iter <= maxIter) { + mgeoms[id].translations[0] = mgeoms[id].translations[0] + + ((mgeoms[id].translations[0] - mgeoms[id].translations[1]) / (float)maxIter); + mgeoms[id].transforms[0] = utilityCore::buildTransformationMatrix(mgeoms[id].translations[0], + mgeoms[id].rotations[0], mgeoms[id].scales[0]); + mgeoms[id].inverseTransforms[0] = glm::inverse(mgeoms[id].transforms[0]); + } +} + +/** +* Scatter a ray with some probabilities according to the material properties. +* For example, a diffuse surface scatters in a cosine-weighted hemisphere. +* A perfect specular surface scatters in the reflected ray direction. +* In order to apply multiple effects to one surface, probabilistically choose +* between them. +* +* The visual effect you want is to straight-up add the diffuse and specular +* components. You can do this in a few ways. This logic also applies to +* combining other types of materias (such as refractive). +* - (NOT RECOMMENDED - converges slowly or badly especially for pure-diffuse +* or pure-specular. In principle this correct, though.) +* Always take a 50/50 split between a diffuse bounce and a specular bounce, +* but multiply the result of either one by 1/0.5 to cancel the 0.5 chance +* of it happening. +* - Pick the split based on the intensity of each color, and multiply each +* branch result by the inverse of that branch's probability (same as above). +* +* This method applies its changes to the Ray parameter `ray` in place. +* It also modifies the color `color` of the ray in place. +* +* You may need to change the parameter list for your purposes! +*/ __host__ __device__ void scatterRay( Ray &ray, glm::vec3 &color, glm::vec3 intersect, glm::vec3 normal, - glm::vec3 emittedColor, const Material &m, thrust::default_random_engine &rng) { - // TODO: implement this. - // A basic implementation of pure-diffuse shading will just call the - // calculateRandomDirectionInHemisphere defined above. + thrust::uniform_real_distribution u01(0, 1); + + if (m.hasRefractive) { + if (u01(rng) < (1.0f - calculateFresnelReflectionCoefficient(ray.direction, normal, m.indexOfRefraction))) { + float eta = 0.0f; + if (!ray.inside) { + eta = 1.0f / m.indexOfRefraction; // Coming from the air + } + + float angle = 1.0f - glm::pow(eta, 2) * (1.0f - glm::pow(glm::dot(normal, ray.direction), 2)); + if (angle < 0.0f) { + // Angle less than zero, so we reflect + ray.direction = calculateReflectionDirection(ray.direction, normal); + ray.origin = intersect + ray.direction * EPSILON; + ray.inside = false; + } + else { + // Here we do a refraction + ray.direction = calculateRefractionDirection(ray.direction, normal, angle, eta); + ray.origin = intersect + ray.direction * 0.001f; // For some reason EPSILON is too small and gives bad results. Need to use larger one + ray.inside = true; + } + } + else { + ray.direction = calculateReflectionDirection(ray.direction, normal); + ray.origin = intersect + ray.direction * EPSILON; + ray.inside = false; + } + } + else if (m.hasReflective) { + //First must determine if this is perfectly specular or not + // is this only when there's an exponent? or when the diffuse is zero? + // for now i will go with the exponent being non zero + if (m.specular.exponent != 0) { + // non perfect + + // Calculate intensity values + float specularIntensity = (m.specular.color.x + m.specular.color.y + m.specular.color.z) / 3.0f; + float diffuseIntensity = (m.color.x + m.color.y + m.color.z) / 3.0f; + + float specularProbability = specularIntensity / (diffuseIntensity + specularIntensity); + + if (u01(rng) <= specularProbability) { + //spec + ray.origin = intersect + normal * EPSILON; + ray.direction = calculateReflectionDirection(ray.direction, normal); + color *= m.specular.color * (1.0f / specularProbability); + ray.inside = false; + } + else { + //diffuse + ray.origin = intersect + normal * EPSILON; + ray.direction = calculateRandomDirectionInHemisphere(normal, rng); + color *= m.color * (1.0f / (diffuseIntensity / (diffuseIntensity + specularIntensity))); + ray.inside = false; + } + } + else { + // perfect mirror + ray.origin = intersect + normal * EPSILON; + ray.direction = calculateReflectionDirection(ray.direction, normal); + color *= m.specular.color; + ray.inside = false; + } + } + else { + // diffuse only + ray.origin = intersect + normal * EPSILON; + ray.direction = calculateRandomDirectionInHemisphere(normal, rng); + color *= m.color; + ray.inside = false; + } } diff --git a/src/intersections.h b/src/intersections.h index 5f3613d..987e3cc 100644 --- a/src/intersections.h +++ b/src/intersections.h @@ -6,139 +6,142 @@ #include "utilities.h" /** - * Handy-dandy hash function that provides seeds for random number generation. - */ +* Handy-dandy hash function that provides seeds for random number generation. +*/ __host__ __device__ inline unsigned int utilhash(unsigned int a) { - a = (a + 0x7ed55d16) + (a << 12); - a = (a ^ 0xc761c23c) ^ (a >> 19); - a = (a + 0x165667b1) + (a << 5); - a = (a + 0xd3a2646c) ^ (a << 9); - a = (a + 0xfd7046c5) + (a << 3); - a = (a ^ 0xb55a4f09) ^ (a >> 16); - return a; + a = (a + 0x7ed55d16) + (a << 12); + a = (a ^ 0xc761c23c) ^ (a >> 19); + a = (a + 0x165667b1) + (a << 5); + a = (a + 0xd3a2646c) ^ (a << 9); + a = (a + 0xfd7046c5) + (a << 3); + a = (a ^ 0xb55a4f09) ^ (a >> 16); + return a; } // CHECKITOUT /** - * Compute a point at parameter value `t` on ray `r`. - * Falls slightly short so that it doesn't intersect the object it's hitting. - */ +* Compute a point at parameter value `t` on ray `r`. +* Falls slightly short so that it doesn't intersect the object it's hitting. +*/ __host__ __device__ glm::vec3 getPointOnRay(Ray r, float t) { - return r.origin + (t - .0001f) * glm::normalize(r.direction); + return r.origin + (t - .0001f) * glm::normalize(r.direction); } /** - * Multiplies a mat4 and a vec4 and returns a vec3 clipped from the vec4. - */ +* Multiplies a mat4 and a vec4 and returns a vec3 clipped from the vec4. +*/ __host__ __device__ glm::vec3 multiplyMV(glm::mat4 m, glm::vec4 v) { - return glm::vec3(m * v); + return glm::vec3(m * v); } // CHECKITOUT /** - * Test intersection between a ray and a transformed cube. Untransformed, - * the cube ranges from -0.5 to 0.5 in each axis and is centered at the origin. - * - * @param intersectionPoint Output parameter for point of intersection. - * @param normal Output parameter for surface normal. - * @return Ray parameter `t` value. -1 if no intersection. - */ +* Test intersection between a ray and a transformed cube. Untransformed, +* the cube ranges from -0.5 to 0.5 in each axis and is centered at the origin. +* +* @param intersectionPoint Output parameter for point of intersection. +* @param normal Output parameter for surface normal. +* @param outside Output param for whether the ray came from outside. +* @return Ray parameter `t` value. -1 if no intersection. +*/ __host__ __device__ float boxIntersectionTest(Geom box, Ray r, - glm::vec3& intersectionPoint, glm::vec3& normal) { - Ray q; - q.origin = multiplyMV(box.inverseTransform, glm::vec4(r.origin , 1.0f)); - q.direction = glm::normalize(multiplyMV(box.inverseTransform, glm::vec4(r.direction, 0.0f))); - - bool outside; - float tmin = -1e38f; - float tmax = 1e38f; - glm::vec3 tmin_n; - glm::vec3 tmax_n; - for (int xyz = 0; xyz < 3; ++xyz) { - float qdxyz = q.direction[xyz]; - /*if (glm::abs(qdxyz) > 0.00001f)*/ { - float t1 = (-0.5f - q.origin[xyz]) / qdxyz; - float t2 = (+0.5f - q.origin[xyz]) / qdxyz; - float ta = glm::min(t1, t2); - float tb = glm::max(t1, t2); - glm::vec3 n; - n[xyz] = t2 < t1 ? +1 : -1; - if (ta > 0 && ta > tmin) { - tmin = ta; - tmin_n = n; - } - if (tb < tmax) { - tmax = tb; - tmax_n = n; - } - } - } - - if (tmax >= tmin && tmax > 0) { - outside = true; - if (tmin <= 0) { - tmin = tmax; - tmin_n = tmax_n; - outside = false; - } - intersectionPoint = multiplyMV(box.transform, glm::vec4(getPointOnRay(q, tmin), 1.0f)); - normal = glm::normalize(multiplyMV(box.transform, - glm::vec4(outside ? tmin_n : -tmin_n, 0.0f))); - return glm::length(r.origin - intersectionPoint); - } - return -1; + glm::vec3 &intersectionPoint, glm::vec3 &normal) { + Ray q; + q.origin = multiplyMV(box.inverseTransform, glm::vec4(r.origin, 1.0f)); + q.direction = glm::normalize(multiplyMV(box.inverseTransform, glm::vec4(r.direction, 0.0f))); + bool outside; + + float tmin = -1e38f; + float tmax = 1e38f; + glm::vec3 tmin_n; + glm::vec3 tmax_n; + for (int xyz = 0; xyz < 3; ++xyz) { + float qdxyz = q.direction[xyz]; + /*if (glm::abs(qdxyz) > 0.00001f)*/ { + float t1 = (-0.5f - q.origin[xyz]) / qdxyz; + float t2 = (+0.5f - q.origin[xyz]) / qdxyz; + float ta = glm::min(t1, t2); + float tb = glm::max(t1, t2); + glm::vec3 n; + n[xyz] = t2 < t1 ? +1 : -1; + if (ta > 0 && ta > tmin) { + tmin = ta; + tmin_n = n; + } + if (tb < tmax) { + tmax = tb; + tmax_n = n; + } + } + } + + if (tmax >= tmin && tmax > 0) { + outside = true; + if (tmin <= 0) { + tmin = tmax; + tmin_n = tmax_n; + outside = false; + } + intersectionPoint = multiplyMV(box.transform, glm::vec4(getPointOnRay(q, tmin), 1.0f)); + normal = glm::normalize(multiplyMV(box.transform, glm::vec4(tmin_n, 0.0f))); + return glm::length(r.origin - intersectionPoint); + } + return -1; } // CHECKITOUT /** - * Test intersection between a ray and a transformed sphere. Untransformed, - * the sphere always has radius 0.5 and is centered at the origin. - * - * @param intersectionPoint Output parameter for point of intersection. - * @param normal Output parameter for surface normal. - * @return Ray parameter `t` value. -1 if no intersection. - */ +* Test intersection between a ray and a transformed sphere. Untransformed, +* the sphere always has radius 0.5 and is centered at the origin. +* +* @param intersectionPoint Output parameter for point of intersection. +* @param normal Output parameter for surface normal. +* @param outside Output param for whether the ray came from outside. +* @return Ray parameter `t` value. -1 if no intersection. +*/ __host__ __device__ float sphereIntersectionTest(Geom sphere, Ray r, - glm::vec3& intersectionPoint, glm::vec3& normal) { - bool outside = false; - float radius = .5; - - glm::vec3 ro = multiplyMV(sphere.inverseTransform, glm::vec4(r.origin, 1.0f)); - glm::vec3 rd = glm::normalize(multiplyMV(sphere.inverseTransform, glm::vec4(r.direction, 0.0f))); - - Ray rt; - rt.origin = ro; - rt.direction = rd; - - float vDotDirection = glm::dot(rt.origin, rt.direction); - float radicand = vDotDirection * vDotDirection - (glm::dot(rt.origin, rt.origin) - pow(radius, 2)); - if (radicand < 0) { - return -1; - } - - float squareRoot = sqrt(radicand); - float firstTerm = -vDotDirection; - float t1 = firstTerm + squareRoot; - float t2 = firstTerm - squareRoot; - - float t = 0; - if (t1 < 0 && t2 < 0) { - return -1; - } else if (t1 > 0 && t2 > 0) { - t = min(t1, t2); - outside = true; - } else { - t = max(t1, t2); - outside = false; - } - - glm::vec3 objspaceIntersection = getPointOnRay(rt, t); - - intersectionPoint = multiplyMV(sphere.transform, glm::vec4(objspaceIntersection, 1.f)); - normal = glm::normalize(multiplyMV(sphere.invTranspose, glm::vec4(objspaceIntersection, 0.f))); - if (!outside) { - normal = -normal; - } - - return glm::length(r.origin - intersectionPoint); -} + glm::vec3 &intersectionPoint, glm::vec3 &normal) { + float radius = .5; + bool outside; + + glm::vec3 ro = multiplyMV(sphere.inverseTransform, glm::vec4(r.origin, 1.0f)); + glm::vec3 rd = glm::normalize(multiplyMV(sphere.inverseTransform, glm::vec4(r.direction, 0.0f))); + + Ray rt; + rt.origin = ro; + rt.direction = rd; + + float vDotDirection = glm::dot(rt.origin, rt.direction); + float radicand = vDotDirection * vDotDirection - (glm::dot(rt.origin, rt.origin) - pow(radius, 2)); + if (radicand < 0) { + return -1; + } + + float squareRoot = sqrt(radicand); + float firstTerm = -vDotDirection; + float t1 = firstTerm + squareRoot; + float t2 = firstTerm - squareRoot; + + float t = 0; + if (t1 < 0 && t2 < 0) { + return -1; + } + else if (t1 > 0 && t2 > 0) { + t = min(t1, t2); + outside = true; + } + else { + t = max(t1, t2); + outside = false; + } + + glm::vec3 objspaceIntersection = getPointOnRay(rt, t); + + intersectionPoint = multiplyMV(sphere.transform, glm::vec4(objspaceIntersection, 1.f)); + normal = glm::normalize(multiplyMV(sphere.invTranspose, glm::vec4(objspaceIntersection, 0.f))); + if (!outside) { + normal = -normal; + } + + return glm::length(r.origin - intersectionPoint); +} \ No newline at end of file diff --git a/src/main.cpp b/src/main.cpp index 77671f4..c9af686 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -6,6 +6,10 @@ static std::string startTimeString; static bool camchanged = false; static float theta = 0, phi = 0; static glm::vec3 cammove; +static bool blur; +static bool dof; +static float focalDistance; +static float apertureRadius; Scene *scene; RenderState *renderState; @@ -37,6 +41,9 @@ int main(int argc, char** argv) { width = renderState->camera.resolution.x; height = renderState->camera.resolution.y; + blur = renderState->camera.blur; + dof = renderState->camera.dof; + // Initialize CUDA and GL components init(); @@ -82,6 +89,12 @@ void runCuda() { cam.position += cammove.x * r + cammove.y * u + cammove.z * v; theta = phi = 0; cammove = glm::vec3(); + cam.blur = blur; + cam.dof = dof; + cam.focalDistance += focalDistance; + focalDistance = 0.0f; + cam.apertureRadius += apertureRadius; + apertureRadius = 0.0f; camchanged = false; } @@ -100,7 +113,7 @@ void runCuda() { // execute the kernel int frame = 0; - pathtrace(pbo_dptr, frame, iteration); + pathtrace(pbo_dptr, frame, iteration, renderState->iterations); // unmap buffer object cudaGLUnmapBufferObject(pbo); @@ -132,6 +145,17 @@ void keyCallback(GLFWwindow* window, int key, int scancode, int action, int mods case GLFW_KEY_S: camchanged = true; cammove -= glm::vec3(0, 0, .1f); break; case GLFW_KEY_R: camchanged = true; cammove += glm::vec3(0, .1f, 0); break; case GLFW_KEY_F: camchanged = true; cammove -= glm::vec3(0, .1f, 0); break; + + // Depth of field controls + case GLFW_KEY_0: camchanged = true; dof = !dof; break; + case GLFW_KEY_1: camchanged = true; focalDistance = -0.1; break; + case GLFW_KEY_2: camchanged = true; focalDistance = +0.1; break; + case GLFW_KEY_3: camchanged = true; apertureRadius = -0.01; break; + case GLFW_KEY_4: camchanged = true; apertureRadius = +0.01; break; + + // Motion blur + // Note: Can only activate and deactivate blur. Requires object parameters to support blur in config files + case GLFW_KEY_B: camchanged = true; blur = !blur; break; } } } diff --git a/src/pathtrace.cu b/src/pathtrace.cu index 4569e2f..18a6f65 100644 --- a/src/pathtrace.cu +++ b/src/pathtrace.cu @@ -1,9 +1,7 @@ #include #include #include -#include #include -#include #include "sceneStructs.h" #include "scene.h" @@ -13,143 +11,307 @@ #include "pathtrace.h" #include "intersections.h" #include "interactions.h" +#include "../stream_compaction/efficient.h" #define FILENAME (strrchr(__FILE__, '/') ? strrchr(__FILE__, '/') + 1 : __FILE__) #define checkCUDAError(msg) checkCUDAErrorFn(msg, FILENAME, __LINE__) void checkCUDAErrorFn(const char *msg, const char *file, int line) { - cudaError_t err = cudaGetLastError(); - if (cudaSuccess == err) { - return; - } - - fprintf(stderr, "CUDA error"); - if (file) { - fprintf(stderr, " (%s:%d)", file, line); - } - fprintf(stderr, ": %s: %s\n", msg, cudaGetErrorString(err)); - exit(EXIT_FAILURE); + cudaError_t err = cudaGetLastError(); + if (cudaSuccess == err) { + return; + } + + fprintf(stderr, "CUDA error"); + if (file) { + fprintf(stderr, " (%s:%d)", file, line); + } + fprintf(stderr, ": %s: %s\n", msg, cudaGetErrorString(err)); + exit(EXIT_FAILURE); } __host__ __device__ thrust::default_random_engine random_engine( - int iter, int index = 0, int depth = 0) { - return thrust::default_random_engine(utilhash((index + 1) * iter) ^ utilhash(depth)); + int iter, int index = 0, int depth = 0) { + int h = utilhash((1 << 31) | (depth << 20) | iter) ^ utilhash(index); + return thrust::default_random_engine(h); } //Kernel that writes the image to the OpenGL PBO directly. __global__ void sendImageToPBO(uchar4* pbo, glm::ivec2 resolution, - int iter, glm::vec3* image) { - int x = (blockIdx.x * blockDim.x) + threadIdx.x; - int y = (blockIdx.y * blockDim.y) + threadIdx.y; - - if (x < resolution.x && y < resolution.y) { - int index = x + (y * resolution.x); - glm::vec3 pix = image[index]; - - glm::ivec3 color; - color.x = glm::clamp((int) (pix.x / iter * 255.0), 0, 255); - color.y = glm::clamp((int) (pix.y / iter * 255.0), 0, 255); - color.z = glm::clamp((int) (pix.z / iter * 255.0), 0, 255); - - // Each thread writes one pixel location in the texture (textel) - pbo[index].w = 0; - pbo[index].x = color.x; - pbo[index].y = color.y; - pbo[index].z = color.z; - } + int iter, glm::vec3* image) { + int x = (blockIdx.x * blockDim.x) + threadIdx.x; + int y = (blockIdx.y * blockDim.y) + threadIdx.y; + + if (x < resolution.x && y < resolution.y) { + int index = x + (y * resolution.x); + glm::vec3 pix = image[index]; + + glm::ivec3 color; + color.x = glm::clamp((int)(pix.x / iter * 255.0), 0, 255); + color.y = glm::clamp((int)(pix.y / iter * 255.0), 0, 255); + color.z = glm::clamp((int)(pix.z / iter * 255.0), 0, 255); + + // Each thread writes one pixel location in the texture (textel) + pbo[index].w = 0; + pbo[index].x = color.x; + pbo[index].y = color.y; + pbo[index].z = color.z; + } } -static Scene *hst_scene; -static glm::vec3 *dev_image; -// TODO: static variables for device memory, scene/camera info, etc -// ... +static Scene *hst_scene = NULL; +static glm::vec3 *dev_image = NULL; +static Ray* dev_rays = NULL; +static Ray* dev_compactionOutput = NULL; +static Geom* dev_geoms = NULL; +static MovingGeom* hst_mgeoms = NULL; +static MovingGeom* dev_mgeoms = NULL; +static Material* dev_materials = NULL; void pathtraceInit(Scene *scene) { - hst_scene = scene; - const Camera &cam = hst_scene->state.camera; - const int pixelcount = cam.resolution.x * cam.resolution.y; + hst_scene = scene; + const Camera &cam = hst_scene->state.camera; + + hst_mgeoms = &(hst_scene->mgeoms)[0]; + + const Material *materials = &(hst_scene->materials)[0]; + const int pixelcount = cam.resolution.x * cam.resolution.y; - cudaMalloc(&dev_image, pixelcount * sizeof(glm::vec3)); - cudaMemset(dev_image, 0, pixelcount * sizeof(glm::vec3)); - // TODO: initialize the above static variables added above + cudaMalloc(&dev_image, pixelcount * sizeof(glm::vec3)); + cudaMemset(dev_image, 0, pixelcount * sizeof(glm::vec3)); - checkCUDAError("pathtraceInit"); + cudaMalloc(&dev_rays, pixelcount * sizeof(Ray)); + cudaMemset(dev_rays, 0, pixelcount * sizeof(Ray)); + + cudaMalloc(&dev_compactionOutput, pixelcount * sizeof(Ray)); + cudaMemset(dev_compactionOutput, 0, pixelcount * sizeof(Ray)); + + cudaMalloc(&dev_materials, pixelcount * sizeof(Material)); + cudaMemcpy(dev_materials, materials, hst_scene->materials.size() * sizeof(Material), cudaMemcpyHostToDevice); + + checkCUDAError("pathtraceInit"); } void pathtraceFree() { - cudaFree(dev_image); - // TODO: clean up the above static variables + cudaFree(dev_image); // no-op if dev_image is null + cudaFree(dev_rays); + cudaFree(dev_compactionOutput); + cudaFree(dev_geoms); + cudaFree(dev_mgeoms); + cudaFree(dev_materials); + + checkCUDAError("pathtraceFree"); +} + +/** + * To accomodate motion blur, we have to load the MotionGeom data to the Geom data on each iteration. + * TODO: Make changes that avoid using this to reduce memory overhead. + */ +Geom *LoadGeoms(MovingGeom *mgeoms, int frame, int numberOfObjects) { + Geom *geoms = (Geom*)malloc(numberOfObjects * sizeof(Geom)); + for (int i = 0; i < numberOfObjects; i++) { + geoms[i].type = hst_mgeoms[i].type; + geoms[i].materialid = hst_mgeoms[i].materialid; + geoms[i].translation = hst_mgeoms[i].translations[frame]; + geoms[i].rotation = hst_mgeoms[i].rotations[frame]; + geoms[i].scale = hst_mgeoms[i].scales[frame]; + geoms[i].transform = hst_mgeoms[i].transforms[frame]; + geoms[i].inverseTransform = hst_mgeoms[i].inverseTransforms[frame]; + geoms[i].invTranspose = hst_mgeoms[i].inverseTransposes[frame]; + } + + cudaMalloc((void**)&dev_geoms, numberOfObjects * sizeof(Geom)); + cudaMemcpy(dev_geoms, geoms, numberOfObjects * sizeof(Geom), cudaMemcpyHostToDevice); - checkCUDAError("pathtraceFree"); + return geoms; } /** - * Example function to generate static and test the CUDA-GL interop. - * Delete this once you're done looking at it! + * Creates a ray through each pixel on the screen. + * Depth of Field: http://mzshehzanayub.blogspot.com/2012/10/gpu-path-tracer.html */ -__global__ void generateStaticDeleteMe(Camera cam, int iter, glm::vec3 *image) { - int x = (blockIdx.x * blockDim.x) + threadIdx.x; - int y = (blockIdx.y * blockDim.y) + threadIdx.y; - - if (x < cam.resolution.x && y < cam.resolution.y) { - int index = x + (y * cam.resolution.x); - - thrust::default_random_engine rng = random_engine(iter, index, 0); - thrust::uniform_real_distribution u01(0, 1); - - // CHECKITOUT: Note that on every iteration, noise gets added onto - // the image (not replaced). As a result, the image smooths out over - // time, since the output image is the contents of this array divided - // by the number of iterations. - // - // Your renderer will do the same thing, and, over time, it will become - // smoother. - image[index] += glm::vec3(u01(rng)); - } +__global__ void InitializeRays(Camera cam, int iter, Ray* rays) { + int x = (blockIdx.x * blockDim.x) + threadIdx.x; + int y = (blockIdx.y * blockDim.y) + threadIdx.y; + int index = x + (y * cam.resolution.x); // index in ray array + + thrust::default_random_engine rng = random_engine(iter, index, 0); + thrust::uniform_int_distribution u01(0.0f, 1.0f); + thrust::uniform_int_distribution uHalf(-0.5f, 0.5f); + + if (cam.dof) { + // Depth of field + glm::vec3 horizontal, middle, vertical; + glm::vec3 pointOnUnitImagePlane, pointOnTrueImagePlane; + + // Compute point on image plane, then plane at focal distance + horizontal = glm::cross(cam.view, cam.up) * glm::sin(-cam.fov.x); + vertical = glm::cross(glm::cross(cam.view, cam.up), cam.view) * glm::sin(-cam.fov.y); + middle = cam.position + cam.view; + + pointOnUnitImagePlane = middle + (((2.0f * ((uHalf(rng) + x) / (cam.resolution.x - 1))) - 1.0f) + * horizontal) + (((2.0f * ((uHalf(rng) + y) / (cam.resolution.y - 1))) - 1.0f) * vertical); + pointOnTrueImagePlane = cam.position + ((pointOnUnitImagePlane - cam.position) * cam.focalDistance); + + // Sample a random point on the lense + float angle = TWO_PI * u01(rng); + float distance = cam.apertureRadius * glm::sqrt(u01(rng)); + glm::vec2 aperture(glm::cos(angle) * distance, glm::sin(angle) * distance); + + rays[index].origin = cam.position + (aperture.x * glm::cross(cam.view, cam.up) + (aperture.y * glm::cross(glm::cross(cam.view, cam.up), cam.view)));; + rays[index].direction = glm::normalize(pointOnTrueImagePlane - rays[index].origin); + } + else { + //No depth of field + float halfResX, halfResY; + + halfResX = cam.resolution.x / 2.0f; + halfResY = cam.resolution.y / 2.0f; + + rays[index].origin = cam.position; + rays[index].direction = cam.view + ((-(halfResX - x + uHalf(rng)) + * sin(cam.fov.x)) / halfResX) * glm::cross(cam.up, cam.view) + (((halfResY - y + uHalf(rng)) * sin(cam.fov.y)) / halfResY) * cam.up; + } + + rays[index].color = glm::vec3(1.0f); + rays[index].pixel_index = index; + rays[index].alive = true; + rays[index].inside = false; +} + +/** + * Traces an individual array for one bounce. + */ +__global__ void TraceBounce(int iter, int depth, glm::vec3 *image, Ray *rays, const Geom *geoms, const int numberOfObjects, const Material *materials) { + // Thread index corresponds to the ray, pixel index is saved member of the ray + int index = blockIdx.x * blockDim.x * blockDim.y + threadIdx.y * blockDim.x + threadIdx.x; + int pixelIndex = rays[index].pixel_index, minGeomIndex = -1; + float t = -1.0f, minT = FLT_MAX; + glm::vec3 minNormal, minIntersectionPoint; + + for (int i = 0; i < numberOfObjects; i++) { + glm::vec3 normal, intersectionPoint; + + if (geoms[i].type == CUBE) { + t = boxIntersectionTest(geoms[i], rays[index], intersectionPoint, normal); + } + else if (geoms[i].type == SPHERE) { + t = sphereIntersectionTest(geoms[i], rays[index], intersectionPoint, normal); + } + else { + printf("Invalid geometry."); + continue; + } + + // Find the closest intersection + if (t != -1.0f && minT > t) { + minT = t; + minNormal = normal; + minIntersectionPoint = intersectionPoint; + minGeomIndex = i; + } + } + + if (minGeomIndex == -1) { + // Nothing was hit + rays[index].alive = false; + image[pixelIndex] += glm::vec3(0.0f); + } + else { + int materialIndex = geoms[minGeomIndex].materialid; + + // Either we hit a light, or we scatter again + if (materials[materialIndex].emittance > EPSILON) { + rays[index].alive = false; + image[pixelIndex] += rays[index].color * materials[materialIndex].color * materials[materialIndex].emittance; + } + else { + thrust::default_random_engine rng = random_engine(iter, pixelIndex, depth); + scatterRay(rays[index], rays[index].color, minIntersectionPoint, minNormal, materials[materialIndex], rng); + } + } } /** * 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) { - const int traceDepth = hst_scene->state.traceDepth; - const Camera &cam = hst_scene->state.camera; - const int pixelcount = cam.resolution.x * cam.resolution.y; - - const int blockSideLength = 8; - const dim3 blockSize(blockSideLength, blockSideLength); - const dim3 blocksPerGrid( - (cam.resolution.x + blockSize.x - 1) / blockSize.x, - (cam.resolution.y + blockSize.y - 1) / blockSize.y); - - /////////////////////////////////////////////////////////////////////////// - - // Recap: - // * Initialize array of path rays (using rays that come out of the camera) - // * You can pass the Camera object to that kernel. - // * For each depth: - // * Compute one ray along each path - many will terminate. - // You'll have to decide how to represent your path rays and how - // you'll mark terminated rays. - // * Add all of the terminated rays' results into the appropriate pixels. - // * Stream compact away all of the terminated paths. - // You may use your implementation or `thrust::remove_if` or its - // cousins. - // * 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 - - generateStaticDeleteMe<<>>(cam, iter, dev_image); - - /////////////////////////////////////////////////////////////////////////// - - // Send results to OpenGL buffer for rendering - sendImageToPBO<<>>(pbo, cam.resolution, iter, dev_image); - - // Retrieve image from GPU - cudaMemcpy(hst_scene->state.image.data(), dev_image, - pixelcount * sizeof(glm::vec3), cudaMemcpyDeviceToHost); - - checkCUDAError("pathtrace"); +void pathtrace(uchar4 *pbo, int frame, int iter, int maxIter) { + const int traceDepth = hst_scene->state.traceDepth; + const Camera &cam = hst_scene->state.camera; + const int numberOfObjects = hst_scene->mgeoms.size(); + const int pixelcount = cam.resolution.x * cam.resolution.y; + + const int blockSideLength = 8; + const int blockSideLengthSquare = pow(blockSideLength, 2); + const dim3 blockSize(blockSideLength, blockSideLength); + const dim3 blocksPerGrid( + (cam.resolution.x + blockSize.x - 1) / blockSize.x, + (cam.resolution.y + blockSize.y - 1) / blockSize.y); + + /////////////////////////////////////////////////////////////////////////// + + if (iter == 1) { + // If the scene has reset, then reset objects in motion to original positions + for (int i = 0; i < numberOfObjects; i++) { + hst_mgeoms[i].translations[0] = hst_mgeoms[i].translations[2]; + } + } + + // Motion blur + Geom *geoms; + if (cam.blur) { + geoms = LoadGeoms(hst_mgeoms, frame, numberOfObjects); + + for (int i = 0; i < numberOfObjects; i++) { + if (hst_mgeoms[i].motionBlur) { + motionBlur(hst_mgeoms, hst_mgeoms[i].id, iter, maxIter); + } + } + } + else { + geoms = LoadGeoms(hst_mgeoms, 2, numberOfObjects); + } + + InitializeRays<<>>(cam, iter, dev_rays); + checkCUDAError("InitializeRays"); + + int currentDepth = 0, rayCount = pixelcount; + while (rayCount > 0 && currentDepth < traceDepth) { + dim3 thread_blocksPerGrid = (rayCount + blockSideLengthSquare - 1) / blockSideLengthSquare; + /*cudaEvent_t start, stop; + float ms_time = 0.0f; + + cudaEventCreate(&start); + cudaEventCreate(&stop);*/ + + //cudaEventRecord(start); + TraceBounce<<>>(iter, currentDepth, dev_image, dev_rays, dev_geoms, numberOfObjects, dev_materials); + //cudaEventRecord(stop); + //cudaEventSynchronize(stop); + //cudaEventElapsedTime(&ms_time, start, stop); + checkCUDAError("TraceBounce"); + + rayCount = StreamCompaction::Efficient::Compact(rayCount, dev_compactionOutput, dev_rays); + /*if (iter == 1) { + printf("Depth %d: Trace took %.5fms and %d out of %d rays remain.\n", currentDepth, ms_time, rayCount, pixelcount); + }*/ + + cudaMemcpy(dev_rays, dev_compactionOutput, rayCount * sizeof(Ray), cudaMemcpyDeviceToDevice); + currentDepth++; + } + + /////////////////////////////////////////////////////////////////////////// + + // Send results to OpenGL buffer for rendering + sendImageToPBO<<>>(pbo, cam.resolution, iter, dev_image); + checkCUDAError("sendImageToPBO"); + + // Retrieve image from GPU + cudaMemcpy(hst_scene->state.image.data(), dev_image, + pixelcount * sizeof(glm::vec3), cudaMemcpyDeviceToHost); + checkCUDAError("cudaMemcpy"); + + // Free geoms here because we are going to keep allocating it on each iteration atm + free(geoms); } + diff --git a/src/pathtrace.h b/src/pathtrace.h index 1241227..9f2cbba 100644 --- a/src/pathtrace.h +++ b/src/pathtrace.h @@ -5,4 +5,4 @@ void pathtraceInit(Scene *scene); void pathtraceFree(); -void pathtrace(uchar4 *pbo, int frame, int iteration); +void pathtrace(uchar4 *pbo, int frame, int iteration, int maxIteration); diff --git a/src/preview.cpp b/src/preview.cpp index 943456f..c489c17 100644 --- a/src/preview.cpp +++ b/src/preview.cpp @@ -139,8 +139,6 @@ bool init() { exit(EXIT_FAILURE); } - width = 800; - height = 800; window = glfwCreateWindow(width, height, "CIS 565 Path Tracer", NULL, NULL); if (!window) { glfwTerminate(); diff --git a/src/scene.cpp b/src/scene.cpp index 0b48f70..c1f3193 100644 --- a/src/scene.cpp +++ b/src/scene.cpp @@ -5,192 +5,257 @@ #include Scene::Scene(string filename) { - cout << "Reading scene from " << filename << " ..." << endl; - cout << " " << endl; - char* fname = (char*)filename.c_str(); - fp_in.open(fname); - if (!fp_in.is_open()) { - cout << "Error reading from file - aborting!" << endl; - throw; - } - while (fp_in.good()) { - string line; - utilityCore::safeGetline(fp_in, line); - if (!line.empty()) { - vector tokens = utilityCore::tokenizeString(line); - if (strcmp(tokens[0].c_str(), "MATERIAL") == 0) { - loadMaterial(tokens[1]); - cout << " " << endl; - } else if (strcmp(tokens[0].c_str(), "OBJECT") == 0) { - loadGeom(tokens[1]); - cout << " " << endl; - } else if (strcmp(tokens[0].c_str(), "CAMERA") == 0) { - loadCamera(); - cout << " " << endl; - } - } - } + cout << "Reading scene from " << filename << " ..." << endl; + cout << " " << endl; + char* fname = (char*)filename.c_str(); + fp_in.open(fname); + if (!fp_in.is_open()) { + cout << "Error reading from file - aborting!" << endl; + throw; + } + while (fp_in.good()) { + string line; + utilityCore::safeGetline(fp_in, line); + if (!line.empty()) { + vector tokens = utilityCore::tokenizeString(line); + if (strcmp(tokens[0].c_str(), "MATERIAL") == 0) { + loadMaterial(tokens[1]); + cout << " " << endl; + } + else if (strcmp(tokens[0].c_str(), "OBJECT") == 0) { + loadGeom(tokens[1]); + cout << " " << endl; + } + else if (strcmp(tokens[0].c_str(), "CAMERA") == 0) { + loadCamera(); + cout << " " << endl; + } + } + } } int Scene::loadGeom(string objectid) { - int id = atoi(objectid.c_str()); - if (id != geoms.size()) { - cout << "ERROR: OBJECT ID does not match expected number of geoms" << endl; - return -1; - } else { - cout << "Loading Geom " << id << "..." << endl; - Geom newGeom; - string line; - - //load object type - utilityCore::safeGetline(fp_in, line); - if (!line.empty() && fp_in.good()) { - if (strcmp(line.c_str(), "sphere") == 0) { - cout << "Creating new sphere..." << endl; - newGeom.type = SPHERE; - } else if (strcmp(line.c_str(), "cube") == 0) { - cout << "Creating new cube..." << endl; - newGeom.type = CUBE; - } - } - - //link material - utilityCore::safeGetline(fp_in, line); - if (!line.empty() && fp_in.good()) { - vector tokens = utilityCore::tokenizeString(line); - newGeom.materialid = atoi(tokens[1].c_str()); - cout << "Connecting Geom " << objectid << " to Material " << newGeom.materialid << "..." << endl; - } - - //load transformations - utilityCore::safeGetline(fp_in, line); - while (!line.empty() && fp_in.good()) { - vector tokens = utilityCore::tokenizeString(line); - - //load tranformations - for (int i = 0; i < 3; i++) { - glm::vec3 translation; - glm::vec3 rotation; - glm::vec3 scale; - utilityCore::safeGetline(fp_in, line); - tokens = utilityCore::tokenizeString(line); - if (strcmp(tokens[0].c_str(), "TRANS") == 0) { - newGeom.translation = glm::vec3(atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str())); - } else if (strcmp(tokens[0].c_str(), "ROTAT") == 0) { - 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())); - } - } - - utilityCore::safeGetline(fp_in, line); - } - - newGeom.transform = utilityCore::buildTransformationMatrix( - newGeom.translation, newGeom.rotation, newGeom.scale); - newGeom.inverseTransform = glm::inverse(newGeom.transform); - newGeom.invTranspose = glm::inverseTranspose(newGeom.transform); - - geoms.push_back(newGeom); - return 1; - } + int id = atoi(objectid.c_str()); + if (id != mgeoms.size()) { + cout << "ERROR: OBJECT ID does not match expected number of geoms" << endl; + return -1; + } + else { + cout << "Loading Geom " << id << "..." << endl; + MovingGeom newGeom; // Switching to a MovingGeom for motion blur + newGeom.id = id; + string line; + + //load object type + utilityCore::safeGetline(fp_in, line); + if (!line.empty() && fp_in.good()) { + if (strcmp(line.c_str(), "sphere") == 0) { + cout << "Creating new sphere..." << endl; + newGeom.type = SPHERE; + } + else if (strcmp(line.c_str(), "cube") == 0) { + cout << "Creating new cube..." << endl; + newGeom.type = CUBE; + } + } + + //link material + utilityCore::safeGetline(fp_in, line); + if (!line.empty() && fp_in.good()) { + vector tokens = utilityCore::tokenizeString(line); + newGeom.materialid = atoi(tokens[1].c_str()); + cout << "Connecting Geom " << objectid << " to Material " << newGeom.materialid << "..." << endl; + } + + //load transformations + int numFrames = 0; + utilityCore::safeGetline(fp_in, line); + vector tempTranslations, tempRotations, tempScales; + bool tempBlur; + while (!line.empty() && fp_in.good()) { + vector tokens = utilityCore::tokenizeString(line); + + if (strcmp(tokens[0].c_str(), "blur") == 0) { + if (atoi(tokens[1].c_str()) == 1) { + tempBlur = true; + } + else { + tempBlur = false; + } + } + else if (strcmp(tokens[0].c_str(), "frame") == 0) { + numFrames++; + } + else if (strcmp(tokens[0].c_str(), "TRANS") == 0) { + tempTranslations.push_back(glm::vec3(atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str()))); + } + else if (strcmp(tokens[0].c_str(), "ROTAT") == 0) { + tempRotations.push_back(glm::vec3(atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str()))); + } + else if (strcmp(tokens[0].c_str(), "SCALE") == 0) { + tempScales.push_back(glm::vec3(atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str()))); + } + + utilityCore::safeGetline(fp_in, line); + } + + numFrames++; // Create extra index for storing original info when scene is reset + // Allocate memory for the geom arrays + newGeom.motionBlur = tempBlur; + newGeom.translations = (glm::vec3*)malloc(numFrames * sizeof(glm::vec3)); + newGeom.rotations = (glm::vec3*)malloc(numFrames * sizeof(glm::vec3)); + newGeom.scales = (glm::vec3*)malloc(numFrames * sizeof(glm::vec3)); + newGeom.transforms = (glm::mat4*)malloc(numFrames * sizeof(glm::mat4)); + newGeom.inverseTransforms = (glm::mat4*)malloc(numFrames * sizeof(glm::mat4)); + newGeom.inverseTransposes = (glm::mat4*)malloc(numFrames * sizeof(glm::mat4)); + + // And finally you can fill them for each frame, and add it onto the list of objects + for (int i = 0; i < numFrames - 1; i++) { + newGeom.translations[i] = tempTranslations[i]; + newGeom.rotations[i] = tempRotations[i]; + newGeom.scales[i] = tempScales[i]; + newGeom.transforms[i] = utilityCore::buildTransformationMatrix(tempTranslations[i], tempRotations[i], tempScales[i]); + newGeom.inverseTransforms[i] = glm::inverse(newGeom.transforms[i]); + newGeom.inverseTransposes[i] = glm::inverseTranspose(newGeom.transforms[i]); + } + + // Save the original twice so we have a backup for when the scene is refreshed + newGeom.translations[2] = tempTranslations[0]; + newGeom.rotations[2] = tempRotations[0]; + newGeom.scales[2] = tempScales[0]; + newGeom.transforms[2] = utilityCore::buildTransformationMatrix(tempTranslations[0], tempRotations[0], tempScales[0]); + newGeom.inverseTransforms[2] = glm::inverse(newGeom.transforms[0]); + newGeom.inverseTransposes[2] = glm::inverseTranspose(newGeom.transforms[0]); + + mgeoms.push_back(newGeom); + return 1; + } } int Scene::loadCamera() { - cout << "Loading Camera ..." << endl; - RenderState &state = this->state; - Camera &camera = state.camera; - float fovy; - - //load static properties - for (int i = 0; i < 5; i++) { - string line; - utilityCore::safeGetline(fp_in, line); - vector tokens = utilityCore::tokenizeString(line); - if (strcmp(tokens[0].c_str(), "RES") == 0) { - camera.resolution.x = atoi(tokens[1].c_str()); - camera.resolution.y = atoi(tokens[2].c_str()); - } else if (strcmp(tokens[0].c_str(), "FOVY") == 0) { - fovy = atof(tokens[1].c_str()); - } else if (strcmp(tokens[0].c_str(), "ITERATIONS") == 0) { - state.iterations = atoi(tokens[1].c_str()); - } else if (strcmp(tokens[0].c_str(), "DEPTH") == 0) { - state.traceDepth = atoi(tokens[1].c_str()); - } else if (strcmp(tokens[0].c_str(), "FILE") == 0) { - state.imageName = tokens[1]; - } - } - - string line; - utilityCore::safeGetline(fp_in, line); - while (!line.empty() && fp_in.good()) { - vector tokens = utilityCore::tokenizeString(line); - - //load camera properties - for (int i = 0; i < 3; i++) { - //glm::vec3 translation; glm::vec3 rotation; glm::vec3 scale; - utilityCore::safeGetline(fp_in, line); - tokens = utilityCore::tokenizeString(line); - if (strcmp(tokens[0].c_str(), "EYE") == 0) { - camera.position = glm::vec3(atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str())); - } else if (strcmp(tokens[0].c_str(), "VIEW") == 0) { - camera.view = glm::vec3(atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str())); - } else if (strcmp(tokens[0].c_str(), "UP") == 0) { - camera.up = glm::vec3(atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str())); - } - } - - utilityCore::safeGetline(fp_in, line); - } - - //calculate fov based on resolution - float yscaled = tan(fovy * (PI / 180)); - float xscaled = (yscaled * camera.resolution.x) / camera.resolution.y; - float fovx = (atan(xscaled) * 180) / PI; - camera.fov = glm::vec2(fovx, fovy); - - //set up render camera stuff - int arraylen = camera.resolution.x * camera.resolution.y; - state.image.resize(arraylen); - std::fill(state.image.begin(), state.image.end(), glm::vec3()); - - cout << "Loaded camera!" << endl; - return 1; + cout << "Loading Camera ..." << endl; + RenderState &state = this->state; + Camera &camera = state.camera; + float fovy; + + //load static properties + for (int i = 0; i < 5; i++) { + string line; + utilityCore::safeGetline(fp_in, line); + vector tokens = utilityCore::tokenizeString(line); + if (strcmp(tokens[0].c_str(), "RES") == 0) { + camera.resolution.x = atoi(tokens[1].c_str()); + camera.resolution.y = atoi(tokens[2].c_str()); + } + else if (strcmp(tokens[0].c_str(), "FOVY") == 0) { + fovy = atof(tokens[1].c_str()); + } + else if (strcmp(tokens[0].c_str(), "ITERATIONS") == 0) { + state.iterations = atoi(tokens[1].c_str()); + } + else if (strcmp(tokens[0].c_str(), "DEPTH") == 0) { + state.traceDepth = atoi(tokens[1].c_str()); + } + else if (strcmp(tokens[0].c_str(), "FILE") == 0) { + state.imageName = tokens[1]; + } + } + + string line; + int numFrames = 0; + utilityCore::safeGetline(fp_in, line); + while (!line.empty() && fp_in.good()) { + vector tokens = utilityCore::tokenizeString(line); + if (strcmp(tokens[0].c_str(), "EYE") == 0) { + camera.position = glm::vec3(atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str())); + } + else if (strcmp(tokens[0].c_str(), "VIEW") == 0) { + camera.view = glm::vec3(atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str())); + } + else if (strcmp(tokens[0].c_str(), "UP") == 0) { + camera.up = glm::vec3(atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str())); + } + else if (strcmp(tokens[0].c_str(), "BLUR") == 0) { + if (atoi(tokens[1].c_str()) == 1) { + camera.blur = true; + } + else { + camera.blur = false; + } + } + else if (strcmp(tokens[0].c_str(), "DOF") == 0) { + if (atoi(tokens[1].c_str()) == 1) { + camera.dof = true; + } + else { + camera.dof = false; + } + } + else if (strcmp(tokens[0].c_str(), "FD") == 0) { + camera.focalDistance = atof(tokens[1].c_str()); + } + else if (strcmp(tokens[0].c_str(), "AR") == 0) { + camera.apertureRadius = atof(tokens[1].c_str()); + } + + utilityCore::safeGetline(fp_in, line); + } + + //calculate fov based on resolution + float yscaled = tan(fovy * (PI / 180)); + float xscaled = (yscaled * camera.resolution.x) / camera.resolution.y; + float fovx = (atan(xscaled) * 180) / PI; + camera.fov = glm::vec2(fovx, fovy); + + //set up render camera stuff + int arraylen = camera.resolution.x * camera.resolution.y; + state.image.resize(arraylen); + std::fill(state.image.begin(), state.image.end(), glm::vec3()); + + cout << "Loaded camera!" << endl; + return 1; } int Scene::loadMaterial(string materialid) { - int id = atoi(materialid.c_str()); - if (id != materials.size()) { - cout << "ERROR: MATERIAL ID does not match expected number of materials" << endl; - return -1; - } else { - cout << "Loading Material " << id << "..." << endl; - Material newMaterial; - - //load static properties - for (int i = 0; i < 10; i++) { - string line; - utilityCore::safeGetline(fp_in, line); - vector tokens = utilityCore::tokenizeString(line); - if (strcmp(tokens[0].c_str(), "RGB") == 0) { - glm::vec3 color( atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str()) ); - newMaterial.color = color; - } else if (strcmp(tokens[0].c_str(), "SPECEX") == 0) { - newMaterial.specular.exponent = atof(tokens[1].c_str()); - } else if (strcmp(tokens[0].c_str(), "SPECRGB") == 0) { - glm::vec3 specColor(atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str())); - newMaterial.specular.color = specColor; - } else if (strcmp(tokens[0].c_str(), "REFL") == 0) { - newMaterial.hasReflective = atof(tokens[1].c_str()); - } else if (strcmp(tokens[0].c_str(), "REFR") == 0) { - newMaterial.hasRefractive = atof(tokens[1].c_str()); - } else if (strcmp(tokens[0].c_str(), "REFRIOR") == 0) { - newMaterial.indexOfRefraction = atof(tokens[1].c_str()); - } else if (strcmp(tokens[0].c_str(), "EMITTANCE") == 0) { - newMaterial.emittance = atof(tokens[1].c_str()); - } - } - materials.push_back(newMaterial); - return 1; - } -} + int id = atoi(materialid.c_str()); + if (id != materials.size()) { + cout << "ERROR: MATERIAL ID does not match expected number of materials" << endl; + return -1; + } + else { + cout << "Loading Material " << id << "..." << endl; + Material newMaterial; + + //load static properties + for (int i = 0; i < 7; i++) { + string line; + utilityCore::safeGetline(fp_in, line); + vector tokens = utilityCore::tokenizeString(line); + if (strcmp(tokens[0].c_str(), "RGB") == 0) { + glm::vec3 color(atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str())); + newMaterial.color = color; + } + else if (strcmp(tokens[0].c_str(), "SPECEX") == 0) { + newMaterial.specular.exponent = atof(tokens[1].c_str()); + } + else if (strcmp(tokens[0].c_str(), "SPECRGB") == 0) { + glm::vec3 specColor(atof(tokens[1].c_str()), atof(tokens[2].c_str()), atof(tokens[3].c_str())); + newMaterial.specular.color = specColor; + } + else if (strcmp(tokens[0].c_str(), "REFL") == 0) { + newMaterial.hasReflective = atof(tokens[1].c_str()); + } + else if (strcmp(tokens[0].c_str(), "REFR") == 0) { + newMaterial.hasRefractive = atof(tokens[1].c_str()); + } + else if (strcmp(tokens[0].c_str(), "REFRIOR") == 0) { + newMaterial.indexOfRefraction = atof(tokens[1].c_str()); + } + else if (strcmp(tokens[0].c_str(), "EMITTANCE") == 0) { + newMaterial.emittance = atof(tokens[1].c_str()); + } + } + materials.push_back(newMaterial); + return 1; + } +} \ No newline at end of file diff --git a/src/scene.h b/src/scene.h index f29a917..3f8707b 100644 --- a/src/scene.h +++ b/src/scene.h @@ -20,7 +20,8 @@ class Scene { Scene(string filename); ~Scene(); - std::vector geoms; + //std::vector geoms; + std::vector mgeoms; std::vector materials; RenderState state; }; diff --git a/src/sceneStructs.h b/src/sceneStructs.h index baa2e30..cb860d2 100644 --- a/src/sceneStructs.h +++ b/src/sceneStructs.h @@ -13,6 +13,10 @@ enum GeomType { struct Ray { glm::vec3 origin; glm::vec3 direction; + glm::vec3 color; + int pixel_index; // the index of the pixel in the array + bool alive; //whether or not hte thread is still running + bool inside; // whether the ray is inside an object or not (changes ior) }; struct Geom { @@ -26,6 +30,21 @@ struct Geom { glm::mat4 invTranspose; }; +//Need to have arrays for our moving geometry +struct MovingGeom { + enum GeomType type; + int id; + int materialid; + bool motionBlur; // if motion blur is enabled or not + int frames; // the frames the geom will move over + glm::vec3 *translations; + glm::vec3 *rotations; + glm::vec3 *scales; + glm::mat4 *transforms; + glm::mat4 *inverseTransforms; + glm::mat4 *inverseTransposes; // Not sure how to use exactly +}; + struct Material { glm::vec3 color; struct { @@ -44,6 +63,10 @@ struct Camera { glm::vec3 view; glm::vec3 up; glm::vec2 fov; + bool blur; + bool dof; + float focalDistance; + float apertureRadius; }; struct RenderState { diff --git a/stream_compaction/CMakeLists.txt b/stream_compaction/CMakeLists.txt index ac358c9..f7ea27d 100644 --- a/stream_compaction/CMakeLists.txt +++ b/stream_compaction/CMakeLists.txt @@ -1,4 +1,6 @@ set(SOURCE_FILES + "efficient.h" + "efficient.cu" ) cuda_add_library(stream_compaction diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu new file mode 100644 index 0000000..aba54b8 --- /dev/null +++ b/stream_compaction/efficient.cu @@ -0,0 +1,191 @@ +#include +#include +#include "efficient.h" + +#define blockSize 128 + +namespace StreamCompaction { +namespace Efficient { + +/** + * Check for CUDA errors; print and exit if there was a problem. + */ +void checkCUDAErrorFn(const char *msg, const char *file, int line) { + cudaError_t err = cudaGetLastError(); + if (cudaSuccess == err) { + return; + } + + fprintf(stderr, "CUDA error"); + if (file) { + fprintf(stderr, " (%s:%d)", file, line); + } + fprintf(stderr, ": %s: %s\n", msg, cudaGetErrorString(err)); + exit(EXIT_FAILURE); +} + +/** +* Maps an array to an array of 0s and 1s for stream compaction. Elements +* which map to 0 will be removed, and elements which map to 1 will be kept. +*/ +__global__ void KernMapToBoolean(int n, int *bools, const Ray *idata) { + int index = threadIdx.x + (blockIdx.x * blockDim.x); + bools[index] = !!idata[index].alive; +} + +/** +* Performs scatter on an array. That is, for each element in idata, +* if bools[idx] == 1, it copies idata[idx] to odata[indices[idx]]. +*/ +__global__ void KernScatter(int n, Ray *odata, const Ray *idata, const int *bools, const int *indices) { + int index = (blockIdx.x * blockDim.x) + threadIdx.x; + + if (bools[index] == 1) { + odata[indices[index]] = idata[index]; + } +} + +/** +* Accumulates the new count of threads for a block with the original. +*/ +__global__ void KernGetBlockCount(int n, int *odata, const int *idata1, const int *idata2) { + int index = (blockIdx.x * blockDim.x) + threadIdx.x; + odata[index] = idata1[(index + 1) * blockSize - 1] + idata2[(index + 1) * blockSize - 1]; +} + +/** +* Increments the block count. +*/ +__global__ void KernIncrementBlock(int n, int *data, const int *increments) { + int index = (blockIdx.x * blockDim.x) + threadIdx.x; + data[index] = data[index] + increments[blockIdx.x]; +} + +/* +* Performs work efficient scan on data in a single GPU block using shared memory. +* Based on the GPU Gems 3 Code found here: +* http://http.developer.nvidia.com/GPUGems3/gpugems3_ch39.html +*/ +__global__ void KernScan(int n, int *odata, const int *idata) { + int index = threadIdx.x; + int offset = 1; + extern __shared__ int temp[]; + + // Copy input data to shared memory + temp[2 * index] = idata[2 * index + (blockIdx.x * blockDim.x * 2)]; + temp[2 * index + 1] = idata[2 * index + 1 + (blockIdx.x * blockDim.x * 2)]; + + // Up sweep + for (int d = n >> 1; d > 0; d >>= 1) { + __syncthreads(); + + if (index < d) { + int ai = offset * (2 * index + 1) - 1; + int bi = offset * (2 * index + 2) - 1; + temp[bi] += temp[ai]; + } + offset *= 2; + } + + // Clear the root + if (index == 0) { + temp[n - 1] = 0; + } + + // Down sweep + for (int d = 1; d < n; d *= 2) { + offset >>= 1; + __syncthreads(); + + if (index < d) { + int ai = offset * (2 * index + 1) - 1; + int bi = offset * (2 * index + 2) - 1; + int t = temp[ai]; + temp[ai] = temp[bi]; + temp[bi] += t; + } + } + __syncthreads(); + + // Write to output array + odata[2 * index + (blockIdx.x * blockDim.x * 2)] = temp[2 * index]; + odata[2 * index + 1 + (blockIdx.x * blockDim.x * 2)] = temp[2 * index + 1]; +} + +void Scan(int n, int *odata, int *idata) { + int blocksPerGrid = (n - 1) / blockSize + 1; + int *dev_idata, *dev_odata; // Padded device memory to handle non power of 2 cases + + cudaMalloc((void**)&dev_idata, blocksPerGrid * blockSize * sizeof(int)); + cudaMemset(dev_idata, 0, blocksPerGrid * blockSize * sizeof(int)); + cudaMemcpy(dev_idata, idata, n * sizeof(int), cudaMemcpyDeviceToDevice); + + cudaMalloc((void**)&dev_odata, blocksPerGrid * blockSize * sizeof(int)); + + if (blocksPerGrid == 1) { + KernScan<<<1, blockSize / 2, blockSize * sizeof(int)>>>(blockSize, dev_odata, dev_idata); + checkCUDAError("KernScan"); + + cudaMemcpy(odata, dev_odata, n * sizeof(int), cudaMemcpyDeviceToDevice); + } + else { + int *dev_increments, *dev_scannedIncrements; + + cudaMalloc((void**)&dev_increments, blocksPerGrid * sizeof(int)); + cudaMalloc((void**)&dev_scannedIncrements, blocksPerGrid * sizeof(int)); + + KernScan<<>>(blockSize, dev_odata, dev_idata); + checkCUDAError("KernScan"); + + int tempBlocksPerGrid = (blocksPerGrid - 1) / blockSize + 1; + KernGetBlockCount<<>>(blocksPerGrid, dev_increments, dev_odata, dev_idata); + checkCUDAError("KernGetBlockCount"); + + // Recursive scan call until we can fit on a single block + Scan(blocksPerGrid, dev_scannedIncrements, dev_increments); + + KernIncrementBlock<<>>(blocksPerGrid * blockSize, dev_odata, dev_scannedIncrements); + checkCUDAError("KernIncrementBlock"); + + cudaMemcpy(odata, dev_odata, n * sizeof(int), cudaMemcpyDeviceToDevice); + + cudaFree(dev_increments); + cudaFree(dev_scannedIncrements); + } + + cudaFree(dev_idata); + cudaFree(dev_odata); +} + +int Compact(int n, Ray *odata, Ray *idata) { + int blocksPerGrid = (n - 1) / blockSize + 1; + int rayCount = 0; + int *dev_bools, *dev_scanData; + + cudaMalloc((void**)&dev_bools, n * sizeof(int)); + cudaMalloc((void**)&dev_scanData, blocksPerGrid * blockSize * sizeof(int)); + + // Map input to boolean values + KernMapToBoolean<<>>(n, dev_bools, idata); + checkCUDAError("KernMapToBoolean"); + + // Scan + Scan(n, dev_scanData, dev_bools); + + // Scatter + KernScatter<<>>(n, odata, idata, dev_bools, dev_scanData); + checkCUDAError("KernScatter"); + + // Get number of rays remaining + int lastScanDataElem, lastBoolElem; + cudaMemcpy(&lastScanDataElem, dev_scanData + n - 1, sizeof(int), cudaMemcpyDeviceToHost); + cudaMemcpy(&lastBoolElem, dev_bools + n - 1, sizeof(int), cudaMemcpyDeviceToHost); + rayCount = lastScanDataElem + lastBoolElem; + + cudaFree(dev_bools); + cudaFree(dev_scanData); + + return rayCount; +} +} +} diff --git a/stream_compaction/efficient.h b/stream_compaction/efficient.h new file mode 100644 index 0000000..6ea4807 --- /dev/null +++ b/stream_compaction/efficient.h @@ -0,0 +1,11 @@ +#pragma once +#include "../src/sceneStructs.h" + +#define FILENAME (strrchr(__FILE__, '/') ? strrchr(__FILE__, '/') + 1 : __FILE__) +#define checkCUDAError(msg) checkCUDAErrorFn(msg, FILENAME, __LINE__) + +namespace StreamCompaction { +namespace Efficient { + int Compact(int n, Ray *odata, Ray *idata); +} +}