diff --git a/CMakeLists.txt b/CMakeLists.txt index 76a65a0..db40597 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -67,7 +67,8 @@ if(${CMAKE_SYSTEM_NAME} MATCHES "Darwin") set(CUDA_PROPAGATE_HOST_FLAGS OFF) endif() -#add_subdirectory(stream_compaction) # TODO: uncomment if using your stream compaction +include_directories(.) +add_subdirectory(stream_compaction) # TODO: uncomment if using your stream compaction add_subdirectory(src) cuda_add_executable(${CMAKE_PROJECT_NAME} @@ -77,7 +78,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 # TODO: uncomment if using your stream compaction ${CORELIBS} ) diff --git a/README.md b/README.md index 5485396..841acbd 100644 --- a/README.md +++ b/README.md @@ -3,284 +3,253 @@ 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) - -### (TODO: Your README) - -*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. - -Instructions (delete me) -======================== - -This is due Thursday, September 24 evening at midnight. - -**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. - -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.** - -**Recommendation:** Every image you save should automatically get a different -filename. Don't delete all of them! For the benefit of your README, keep a -bunch of them around so you can pick a few to document your progress at the -end. - -### Contents - -* `src/` C++/CUDA source files. -* `scenes/` Example scene description files. -* `img/` Renders of example scene description files. - (These probably won't match precisely with yours.) -* `external/` Includes and static libraries for 3rd party libraries. - - -### Running the code - -The main function requires a scene description file. Call the program with -one as an argument: `cis565_path_tracer scene/sphere.txt`. -(In Visual Studio, `../scene/sphere.txt`.) - -If you are using Visual Studio, you can set this in the Debugging > Command -Arguments section in the Project properties. Make sure you get the path right - -read the console for errors. - -#### Controls - -* Esc to save an image and exit. -* Space to save an image. Watch the console for the output filename. -* W/A/S/D and R/F move the camera. Arrow keys rotate. - -## Requirements - -**Ask on the mailing list for clarifications.** - -In this project, you are given code for: - -* Loading and reading the scene description format -* Sphere and box intersection functions -* Support for saving images -* Working CUDA-GL interop for previewing your render while it's running -* A function which generates random screen noise (instead of an actual render). - -You will need to implement the following features: - -* Raycasting from the camera into the scene through an imaginary grid of pixels - (the screen) - * Implement antialiasing (by jittering rays within each pixel) -* Diffuse surfaces -* Perfectly specular-reflective (mirrored) surfaces - * See notes on diffuse/specular in `scatterRay` -* Stream compaction optimization. You may use any of: - * Your global-memory work-efficient stream compaction implementation. - * A shared-memory work-efficient stream compaction (see below). - * `thrust::remove_if` or any of the other Thrust stream compaction functions. - -You are also required to implement at least 2 of the following features. -Please ask if you need good references (they will be added to this README -later on). If you find good references, share them! **Extra credit**: implement -more features on top of the 2 required ones, with point value up to +20/100 at -the grader's discretion (based on difficulty and coolness). - -* Work-efficient stream compaction using shared memory across multiple blocks - (See *GPU Gems 3* Chapter 39). -* These 2 smaller features: - * Refraction (e.g. glass/water) with Frensel effects using Schlick's - approximation or more accurate methods - * Physically-based depth-of-field (by jittering rays within an aperture) - * Recommended but not required: non-perfect specular surfaces -* Texture mapping -* Bump mapping -* Direct lighting (by taking a final ray directly to a random point on an - emissive object acting as a light source) -* Some method of defining object motion, and motion blur -* Subsurface scattering -* Arbitrary mesh loading and rendering (e.g. `obj` files). You can find these - online or export them from your favorite 3D modeling application. - With approval, you may use a third-party OBJ loading code to bring the data - into C++. - * You can use the triangle intersection function `glm::intersectRayTriangle`. - -This 'extra features' list is not comprehensive. If you have a particular idea -you would like to implement (e.g. acceleration structures, etc.), please -contact us first. - -For each extra feature, you must provide the following analysis: - -* Overview write-up of the feature -* Performance impact of the feature -* If you did something to accelerate the feature, what did you do and why? -* Compare your GPU version of the feature to a HYPOTHETICAL CPU version - (you don't have to implement it!) Does it benefit or suffer from being - implemented on the GPU? -* How might this feature be optimized beyond your current implementation? - -## Base Code Tour - -You'll be working in the following files. Look for important parts of the code: -search for `CHECKITOUT`. You'll have to implement parts labeled with `TODO`. -(But don't let these constrain you - you have free rein!) - -* `src/pathtrace.cu`: path tracing kernels, device functions, and calling code - * `pathtraceInit` initializes the path tracer state - it should copy - scene data (e.g. geometry, materials) from `Scene`. - * `pathtraceFree` frees memory allocated by `pathtraceInit` - * `pathtrace` performs one iteration of the rendering - it handles kernel - launches, memory copies, transferring some data, etc. - * See comments for a low-level path tracing recap. -* `src/intersections.h`: ray intersection functions - * `boxIntersectionTest` and `sphereIntersectionTest`, which take in a ray and - a geometry object and return various properties of the intersection. -* `src/interactions.h`: ray scattering functions - * `calculateRandomDirectionInHemisphere`: a cosine-weighted random direction - in a hemisphere. Needed for implementing diffuse surfaces. - * `scatterRay`: this function should perform all ray scattering, and will - call `calculateRandomDirectionInHemisphere`. See comments for details. -* `src/main.cpp`: you don't need to do anything here, but you can change the - program to save `.hdr` image files, if you want (for postprocessing). - -### Generating random numbers +* Tongbo Sui +* Tested on: Windows 10, i5-3320M @ 2.60GHz 8GB, NVS 5400M 2GB (Personal) + +## Overview +A GPU based path tracer that supports several common material effects. Stream compaction performance is optimized to loosely match that of Thrust. + +## Features +* **Work-efficient stream compaction using shared memory** + * Work-efficient stream compaction for arbitrary input size. +* **Raycasting** + * Cast rays for each pixel from the camera into the scene through an imaginary grid of pixels + * Casting to image plane at a distance calculated from camera FOV. + +###### Raycasting debug. Left shows raycasting directions mapped from `(x,y,z)` to `(r,g,b)`. Right shows the first intersection normal. +![](img/raycast_debug.png) + +* **Diffuse surfaces** + * Cosine weighted diffuse reflection. +###### Diffuse and perfect specular surfaces. The perfect specular shown in this image is incorrect. It's been replaced by non-perfect specular implementation. +![](img/diff_spec.png) + +* **Non-perfect specular surfaces** + * Cosine weighted reflection restricted by specular exponent. Perfectly specular-reflective (mirrored) surfaces would be a non-perfect surface with very large (toward positive infinity) specular exponent. + * Instead of trying to find a biased function, simply jitter the reflection normal. + * http://www.cs.cornell.edu/courses/cs4620/2012fa/lectures/37raytracing.pdf + * *Performance impact*: about twice as many as computations needed. Execution time almost doubled. + * *If on CPU*: looping through each pixel would have `O(n)` time with some constant `c`. This extra effect will increase `c` by at least 2x. + * *Possible further optimization*: faster method for randomizing bounce direction; a better weight function for combining all possible effects. + +###### Non-perfect specular surfaces. Higher specular exponent yields clearer reflection images. +![](img/specs.png) + +* **Refraction with Frensel effects** + * Refraction based on Index of Refraction (IOR), with reflection intensity approximated using Schlick's approximation. + * https://en.wikipedia.org/wiki/Schlick's_approximation + * *Performance impact*: extra calculations are needed as it calculates all possible effects for the material. + * *If on CPU*: looping through each pixel would have `O(n)` time with some constant `c`. This extra effect increases the constant `c` dramatically, as it calculates both reflection and refraction, and split the iteration among diffuse, SSS, reflection, refraction. + * *Possible further optimization*: faster method for randomizing bounce direction. + +###### Refractions. Higher IOR results in more distorted images. +![](img/ior_comparision.png) + +* **Subsurface scattering (SSS)** + * Propagates light under object skin and send them off in randomized directions with an offset. + +###### Subsurface scattering(SSS). Object "thickness" and distance to light source are two major factors of SSS intensity. +![](img/sss_comparison.png) + + * SSS for reflective material is also implemented. + +###### Subsurface scattering(SSS) with reflections. Materials are combined to produce the end result. +![](img/sss_refl_comparison.png) + + * *Optimization*: + * Implementation is a simplified version of Dipole approximation. + * https://graphics.stanford.edu/papers/bssrdf/bssrdf.pdf + * Actual accurate SSS calculation would involve passing geometry object into the function and performing multiple extra intersection tests. This slows the entire rendering process by a factor of ~3. Major reason is that the geometry object is big in memory, and intersection tests are slow. + * Reduced memory overhead by guessing ray-out position. Therefore intersection tests are not performed within the scatter method. Passing geometry all the way into scatter function is avoided. + * If the guessed position is still inside the object (when it should be outside), it will be catched at the next trace depth, where the intersection point is automatically available. Color is multiplied by a factor to cancel out guessing failures. + * *Performance impact*: like other material effects, it increases color calculation by a constant time due to extra computation on each pixel. Moreover, all extra material effects add execution time for the kernel. Due to the fact that warps can only execute all codes then mask out unwanted results, each additional effect equally add more tasks to each thread. Taking all effects (specular, refraction, SSS) in to account, the execution time increased by about 400% compared to a diffuse-only implementation. + * *If on CPU*: looping through each pixel would have `O(n)` time with some constant `c`. This extra effect increases the constant `c` over the loop. + * *Possible further optimization*: a better weight function rather than split 50/50. + +###### More SSS. Object "thickness" and distance to light source are two major factors of SSS intensity. Below shows results from multiple distances. "Thickness" is more or less conceptually equivalent to unlit side's distance to light source. +![](img/sss_fun.png) + +* **Antialiasing** + * Oversampling at each iteration to smooth out jagged edges in the render. + * *Performance impact*: Effectively increases render time; proportional to # of oversampling passes. At each sampling pass, the jittering kernel takes ~440 microsec to finish. Plus oversampling needs averaging, which takes ~140 microseconds For a 3-pass render the overhead will be ~1460 microseconds per iteration. Furthermore, each pass runs the entire tracing once, therefore the final execution time for a 3-pass render will be at least 3x the normal tracing time plus overhead. + * *If on CPU*: aside from initializing each ray in the beginning, a CPU version would loop again at each pixel to randomize the ray, which is `O(n)`. On a GPU each thread takes care of one pixel, which makes the process run in constant time. + * *Possible further optimization*: same as antialiasing since they use the same routine. + +###### Antialiasing comparison. Notice the jagged refraction on the right. +![](img/antialiasing_comparison.png) + +* **Physically-based depth-of-field** + * Depth-of-field effect using antialiasing routines but with different jittering radius and method. + * *Performance impact*: same as antialiasing since they use the same routine. + * *If on CPU*: same as antialiasing. + * *Possible further optimization*: same as antialiasing. + +###### Depth-of-field comparison. Objects outside the focal plane are blurred. +![](img/dof_comparison.png) + +## Analysis + +* Cache vs. direct access: + * Creating a temporary variable for caching elements in large arrays is not always effective in CUDA kernels. In fact, sometimes directly passing the array element around results in much better performance in terms of both exec time and memory access. + * For example, caching a `(ray, color)` pair for later computation only reduces performance in long computations such as intersection test. On the other hand, caching geometries for repetitive access across different threads increases performance. However, a caching variable increases performance in simple and short kernels like pixel painting. + * The rational is that for long kernels, register resources are often depleted easily. A "cache" variable in this case would further worsen the situation by taking up a lot of registers in the first hand, leaving much less registers for the rest of the computations. On the other hand, long kernels tend to fill up the register limitation, causing much less warps available to run in parallel. Less warps means less chance to hide memory access latency. + * This effectively cancelled out the reduction in memory access. For short kernels, a cache variable sometimes can successfully hide memory access latency by carefully ordering operation orders. + +* Stream compaction, open scenes vs. closed scenes: + * Stream compaction greatly reduces input size in open scenes. + * Open scene renders much faster. More rays got terminated due to rays shooting side ways are off to the ambient and thus no hits. + * Closed scene renders much slower. Less rays got terminated because all rays will hit at least a wall, if not a light source. + * As a result, closed scene is much brighter. + +###### Open vs. closed scene. Closed scene has many more active (unterminated) rays in each depth. +![](img/ray_count.png) + +* Thrust vs. custom work-efficient compaction: + * Thrust is still faster performance-wise. + * Custom implementation only loosely match the performance of Thrust (~5ms slower per trace depth). + * Interestingly, the performance gap doesn't seem to change a lot with varying input size. Perhaps the bottlenecks are some fixed calculations that are sub-optimal. + +###### Stream compaction performance. Time is for a complete trace (intersection test, paint, stream compaction, scatter). Difference is the stream compaction performance difference. +![](img/comparison_new.png) + +## Optimization + +**Baseline: `cornell6`, 200*200, scan block size 64** + +* **Scan** + * Occupacy with block size + * Block size needs to be `2^n`: pick 128 based on graph prediction, occupacy goes from 33.3% to 41.28% + +###### Occupacy as a function of block size. +![](img/scan_occupacy.png) + + * Occupacy with register counts: wasn't able to reduce register count, but reduced execution time. + * Remove shorthand variable for `threadIdx.x`: + * No effect (should have used one less register). + * Store `threadIdx-1` to avoid repetitive calculations: + * No effect on registers + * Reduced execution time from ~7700 to ~4000 microsec. + * Change loop ceiling `ilog2ceil(n)` to a compile-time constant: + * No visible effect, but two less calculations per thread. + * Pre-populate `2^d` values in shared memory: + * No effect on registers + * Reduced execution time from ~4000 to ~330 microsec. + * *Conclusion*: + * Execution time: dropped from ~7700 to ~330 microsec; 23x speed up. + * Trade-off: achieved occupacy dropped from 41.28% to 40.6% + +* **Scan scatter (stream compaction)** + * Changed block size based on profiling tool prediction: + * Occupacy goes from 33.3% to 100% + * Put extra temporary variables to hold array elements to be used later + * Successfully utilized larger number of warps and hid memory access latency + * In issue stall reasons, reduced proportion of memory dependency from 70% to 26.4% + * 1.92x execution time speed up + * Trade-off: occupacy dropped to 83.3% +* **Intersection test** + * Block size left unchanged as it's optimal. + * Pre-cache geometries in shared memory: + * Reduced access to device memory by 30% in data size, without obvious change in execution time. + * Higher L1 cache hitrate as a result. + * In some cases execution becomes faster but the improvement is not consistent. + * LIMITATION: Current code allows only 64 geometries in the shared memory, which is equal to the block size; this is fine for non-mesh geometries and less than 64 geometries. For bigger scenes, extra codes for loading more geometries will be needed for it to work. The extra codes shouldn't be hard, but would have an obvious impact on the performance due to more memory access. + * Move temporary variable declarations out of for-loop; change all parameters of intersection test to pass by reference: + * On top of above memory improvement, further reduced access to device memory by 30% in data size. + * However, this only reduces access to thread's temporary variables. Therefore the reduction might not scale with scene window size. + * 1.15x speed up on execution time (~200 microsec). + * The speed up and memory improvement are cancelled if the methods called by intersection tests also have all of their parameters passed by reference. + * Remove temporary variable that stores `(ray, color)` pair, which is stored as a struct: + * Contrary to the belief that the temporary variable is caching and is faster, direct access is actually better. + * Further reduced device memory access by 40% + * Further reduced execution time by ~150 microsec. + * Reduced register count by 1. + * Increased L1 cache hitrate: from 65% to 85% + * Trade-off: increased global memory replay overhead from 25% to 34.7% +* **Ray scatter** + * Remove temporary variable that stores `(ray, color)` pair: + * Reduced register count: from 63 to 55. + * 1.63x execution time speed up. + * 50% less device memory access (data size). + * Increased L2 cache hitrate: from 31.6% to 40.6% + * Reduced global memory replay overhead: from 49.2% to 39.3% + * Refactor `scatterRay` to remove redundant parameters: + * Minor improvements regarding items above. +* **Path termination** + * Remove temporary variable for `(ray, color)` pair and material: + * Reduced register count: from 29 to 12. + * 5.86x execution time speed up. + * Trade-off: high global memory replay overhead (47.5%) +* **Camera raycasting** + * Remove temporary variable for `(ray, color)` pair: + * 1.5x execution time speed up. +* `nvprof` + * Some thoughts were put on reducing overall execution time of stream compaction by working on the methods outside of kernels. + * More analysis on API calls showed that calls to `cudaMalloc` and `cudaFree` are the most time consuming ones. Removing excessive calls to these APIs in the stream compaction method reduces overall execution time by ~15 ms (first trace depth). + * Overall performance is improved by a considerable amount by simply reducing calls to CUDA memory operations. + +###### Optimization baseline. +![](img/baseline.png) + +###### Optimized result. +![](img/optimized_new.png) + +## Appendix + +### Scenes list +* `cornell1`: mixed objects (specular, refraction, diffuse, caustic) +* `cornell2`: SSS, same size spheres, mixed distances to light +* `cornell3`: SSS + specular +* `cornell4`: SSS, same size cubes, mixed distances to light +* `cornell5`: mixed objects, closed scene, camera inside box +* `cornell6`: single sphere, very small window size, for optimization testing + +### Performance test data ``` -thrust::default_random_engine rng(hash(index)); -thrust::uniform_real_distribution u01(0, 1); -float result = u01(rng); +Open scene +800 * 800 cornell1 +Depth: 0 / Grid size: 625125 +Depth: 1 / Grid size: 497276 +Depth: 2 / Grid size: 406024 +Depth: 3 / Grid size: 326613 +Depth: 4 / Grid size: 264707 +Depth: 5 / Grid size: 214660 +Depth: 6 / Grid size: 174592 +Depth: 7 / Grid size: 141986 ``` - -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); +Closed scene +800 * 800 cornell5 +Depth: 0 / Grid size: 624854 +Depth: 1 / Grid size: 608350 +Depth: 2 / Grid size: 592705 +Depth: 3 / Grid size: 579828 +Depth: 4 / Grid size: 568568 +Depth: 5 / Grid size: 558072 +Depth: 6 / Grid size: 547953 +Depth: 7 / Grid size: 538102 +``` +``` +Custom work-efficient stream compaction +800 * 800 cornell1 +Depth: 0 / Grid size: 625107 / Time: 97.423164 +Depth: 1 / Grid size: 511896 / Time: 99.192924 +Depth: 2 / Grid size: 434527 / Time: 84.930527 +Depth: 3 / Grid size: 359734 / Time: 73.663490 +Depth: 4 / Grid size: 303851 / Time: 62.820320 +Depth: 5 / Grid size: 255227 / Time: 54.307522 +Depth: 6 / Grid size: 216026 / Time: 47.650337 +Depth: 7 / Grid size: 182477 / Time: 41.597088 +``` +``` +Thrust +800 * 800 cornell1 +Depth: 0 / Grid size: 625107 / Time: 92.244354 +Depth: 1 / Grid size: 512883 / Time: 94.210144 +Depth: 2 / Grid size: 435087 / Time: 80.382782 +Depth: 3 / Grid size: 359297 / Time: 67.422592 +Depth: 4 / Grid size: 303702 / Time: 57.761951 +Depth: 5 / Grid size: 255254 / Time: 49.950497 +Depth: 6 / Grid size: 215942 / Time: 43.392319 +Depth: 7 / Grid size: 182385 / Time: 38.022335 ``` - -### Notes on GLM - -This project uses GLM for linear algebra. - -On NVIDIA cards pre-Fermi (pre-DX12), you may have issues with mat4-vec4 -multiplication. If you have one of these cards, be careful! If you have issues, -you might need to grab `cudamat4` and `multiplyMV` from the -[Fall 2014 project](https://github.com/CIS565-Fall-2014/Project3-Pathtracer). -Let us know if you need to do this. - -### Scene File Format - -This project uses a custom scene description format. Scene files are flat text -files that describe all geometry, materials, lights, cameras, and render -settings inside of the scene. Items in the format are delimited by new lines, -and comments can be added using C-style `// comments`. - -Materials are defined in the following fashion: - -* MATERIAL (material ID) //material header -* RGB (float r) (float g) (float b) //diffuse color -* SPECX (float specx) //specular exponent -* SPECRGB (float r) (float g) (float b) //specular color -* REFL (bool refl) //reflectivity flag, 0 for no, 1 for yes -* REFR (bool refr) //refractivity flag, 0 for no, 1 for yes -* REFRIOR (float ior) //index of refraction for Fresnel effects -* EMITTANCE (float emittance) //the emittance of the material. Anything >0 - makes the material a light source. - -Cameras are defined in the following fashion: - -* CAMERA //camera header -* RES (float x) (float y) //resolution -* FOVY (float fovy) //vertical field of view half-angle. the horizonal angle is calculated from this and the reslution -* ITERATIONS (float interations) //how many iterations to refine the image, - only relevant for supersampled antialiasing, depth of field, area lights, and - other distributed raytracing applications -* DEPTH (int depth) //maximum depth (number of times the path will bounce) -* FILE (string filename) //file to output render to upon completion -* EYE (float x) (float y) (float z) //camera's position in worldspace -* VIEW (float x) (float y) (float z) //camera's view direction -* UP (float x) (float y) (float z) //camera's up vector - -Objects are defined in the following fashion: - -* OBJECT (object ID) //object header -* (cube OR sphere OR mesh) //type of object, can be either "cube", "sphere", or - "mesh". Note that cubes and spheres are unit sized and centered at the - origin. -* material (material ID) //material to assign this object -* TRANS (float transx) (float transy) (float transz) //translation -* ROTAT (float rotationx) (float rotationy) (float rotationz) //rotation -* SCALE (float scalex) (float scaley) (float scalez) //scale - -Two examples are provided in the `scenes/` directory: a single emissive sphere, -and a simple cornell box made using cubes for walls and lights and a sphere in -the middle. - -## Third-Party Code Policy - -* Use of any third-party code must be approved by asking on our Google Group. -* If it is approved, all students are welcome to use it. Generally, we approve - use of third-party code that is not a core part of the project. For example, - for the path tracer, we would approve using a third-party library for loading - models, but would not approve copying and pasting a CUDA function for doing - refraction. -* Third-party code **MUST** be credited in README.md. -* Using third-party code without its approval, including using another - student's code, is an academic integrity violation, and will, at minimum, - result in you receiving an F for the semester. - -## README - -Please see: [**TIPS FOR WRITING AN AWESOME README**](https://github.com/pjcozzi/Articles/blob/master/CIS565/GitHubRepo/README.md) - -* Sell your project. -* Assume the reader has a little knowledge of path tracing - don't go into - detail explaining what it is. Focus on your project. -* Don't talk about it like it's an assignment - don't say what is and isn't - "extra" or "extra credit." Talk about what you accomplished. -* Use this to document what you've done. -* *DO NOT* leave the README to the last minute! It is a crucial part of the - project, and we will not be able to grade you without a good README. - -In addition: - -* This is a renderer, so include images that you've made! -* Be sure to back your claims for optimization with numbers and comparisons. -* If you reference any other material, please provide a link to it. -* You wil not be graded on how fast your path tracer runs, but getting close to - real-time is always nice! -* If you have a fast GPU renderer, it is very good to show case this with a - video to show interactivity. If you do so, please include a link! - -### 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/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/aliased.png b/img/aliased.png new file mode 100644 index 0000000..b21d341 Binary files /dev/null and b/img/aliased.png differ diff --git a/img/antialiased.png b/img/antialiased.png new file mode 100644 index 0000000..77fe2b4 Binary files /dev/null and b/img/antialiased.png differ diff --git a/img/antialiasing_comparison.png b/img/antialiasing_comparison.png new file mode 100644 index 0000000..ed929b2 Binary files /dev/null and b/img/antialiasing_comparison.png differ diff --git a/img/baseline.png b/img/baseline.png new file mode 100644 index 0000000..6007815 Binary files /dev/null and b/img/baseline.png differ diff --git a/img/closed_scene.png b/img/closed_scene.png new file mode 100644 index 0000000..1c7dcf5 Binary files /dev/null and b/img/closed_scene.png differ diff --git a/img/comparison.png b/img/comparison.png new file mode 100644 index 0000000..abee6cc Binary files /dev/null and b/img/comparison.png differ diff --git a/img/comparison_new.png b/img/comparison_new.png new file mode 100644 index 0000000..8bbd109 Binary files /dev/null and b/img/comparison_new.png differ diff --git a/img/depth_of_field.png b/img/depth_of_field.png new file mode 100644 index 0000000..34ed80e Binary files /dev/null and b/img/depth_of_field.png differ diff --git a/img/diff_spec.png b/img/diff_spec.png new file mode 100644 index 0000000..28a2965 Binary files /dev/null and b/img/diff_spec.png differ diff --git a/img/diffuse_check.png b/img/diffuse_check.png new file mode 100644 index 0000000..74c4cdf Binary files /dev/null and b/img/diffuse_check.png differ diff --git a/img/dof_comparison.png b/img/dof_comparison.png new file mode 100644 index 0000000..3c28fc3 Binary files /dev/null and b/img/dof_comparison.png differ diff --git a/img/hit_normal.png b/img/hit_normal.png new file mode 100644 index 0000000..40f2e96 Binary files /dev/null and b/img/hit_normal.png differ diff --git a/img/ior_comparision.png b/img/ior_comparision.png new file mode 100644 index 0000000..e01a6b0 Binary files /dev/null and b/img/ior_comparision.png differ diff --git a/img/open_comparison.png b/img/open_comparison.png new file mode 100644 index 0000000..be5fd4e Binary files /dev/null and b/img/open_comparison.png differ diff --git a/img/optimized.png b/img/optimized.png new file mode 100644 index 0000000..513bf9b Binary files /dev/null and b/img/optimized.png differ diff --git a/img/optimized_new.png b/img/optimized_new.png new file mode 100644 index 0000000..98e5d97 Binary files /dev/null and b/img/optimized_new.png differ diff --git a/img/ray_count.png b/img/ray_count.png new file mode 100644 index 0000000..db626c0 Binary files /dev/null and b/img/ray_count.png differ diff --git a/img/raycast_debug.png b/img/raycast_debug.png new file mode 100644 index 0000000..5607679 Binary files /dev/null and b/img/raycast_debug.png differ diff --git a/img/raycasting.png b/img/raycasting.png new file mode 100644 index 0000000..e0ccaff Binary files /dev/null and b/img/raycasting.png differ diff --git a/img/refract_diamond.png b/img/refract_diamond.png new file mode 100644 index 0000000..3967fdd Binary files /dev/null and b/img/refract_diamond.png differ diff --git a/img/refract_ice.png b/img/refract_ice.png new file mode 100644 index 0000000..a78aeaa Binary files /dev/null and b/img/refract_ice.png differ diff --git a/img/scan_occupacy.png b/img/scan_occupacy.png new file mode 100644 index 0000000..6cc8e75 Binary files /dev/null and b/img/scan_occupacy.png differ diff --git a/img/REFERENCE_sphere.5000samp.png b/img/simple_sphere.png similarity index 60% rename from img/REFERENCE_sphere.5000samp.png rename to img/simple_sphere.png index 0628927..96fdd4a 100644 Binary files a/img/REFERENCE_sphere.5000samp.png and b/img/simple_sphere.png differ diff --git a/img/specs.png b/img/specs.png new file mode 100644 index 0000000..1178812 Binary files /dev/null and b/img/specs.png differ diff --git a/img/sss_check.png b/img/sss_check.png new file mode 100644 index 0000000..9bbfa3c Binary files /dev/null and b/img/sss_check.png differ diff --git a/img/sss_comparison.png b/img/sss_comparison.png new file mode 100644 index 0000000..c511ae6 Binary files /dev/null and b/img/sss_comparison.png differ diff --git a/img/sss_distance_cube.png b/img/sss_distance_cube.png new file mode 100644 index 0000000..8ecfe7b Binary files /dev/null and b/img/sss_distance_cube.png differ diff --git a/img/sss_distance_sphere.png b/img/sss_distance_sphere.png new file mode 100644 index 0000000..84c110b Binary files /dev/null and b/img/sss_distance_sphere.png differ diff --git a/img/sss_fun.png b/img/sss_fun.png new file mode 100644 index 0000000..e362c00 Binary files /dev/null and b/img/sss_fun.png differ diff --git a/img/sss_refl_comparison.png b/img/sss_refl_comparison.png new file mode 100644 index 0000000..f6b905b Binary files /dev/null and b/img/sss_refl_comparison.png differ diff --git a/img/sss_reflection.png b/img/sss_reflection.png new file mode 100644 index 0000000..19b2884 Binary files /dev/null and b/img/sss_reflection.png differ diff --git a/scenes/cornell.txt b/scenes/cornell.txt index a18639e..7bfd13b 100644 --- a/scenes/cornell.txt +++ b/scenes/cornell.txt @@ -104,4 +104,4 @@ sphere material 1 TRANS -1 4 -1 ROTAT 0 0 0 -SCALE 3 3 3 +SCALE 3 3 3 \ No newline at end of file diff --git a/scenes/cornell1.txt b/scenes/cornell1.txt new file mode 100644 index 0000000..c383004 --- /dev/null +++ b/scenes/cornell1.txt @@ -0,0 +1,206 @@ +// Emissive material (light) +MATERIAL 0 +RGB 1 1 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 1.5 + +// 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 30 +SPECRGB 1 1 1 +REFL 1 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Diffuse green +MATERIAL 3 +RGB .35 .85 .35 +SPECEX 1000 +SPECRGB 1 1 1 +REFL 1 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Emissive material (light) +MATERIAL 4 +RGB 1 1 0 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 1 + +// Glossy material +MATERIAL 5 +RGB .35 .35 .85 +SPECEX 5 +SPECRGB .8 .8 .8 +REFL 1 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Glossy material +MATERIAL 6 +RGB .35 .85 .35 +SPECEX 0 +SPECRGB .8 .8 .8 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Mirror white +MATERIAL 7 +RGB .98 .98 .98 +SPECEX 100 +SPECRGB 1 1 1 +REFL 1 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Refraction +MATERIAL 8 +RGB .847 .949 .9529 +SPECEX 1000 +SPECRGB 1 1 1 +REFL 0 +REFR 1 +REFRIOR 1.309 +EMITTANCE 0 + +// Camera +CAMERA +RES 800 800 +FOVY 45 +ITERATIONS 5000 +DEPTH 8 +FILE cornell +EYE 0.0 5 14 +VIEW 0 0 -1 +UP 0 1 0 +DOF 9 + + +// Ceiling light +OBJECT 0 +cube +material 0 +TRANS 0 10 0 +ROTAT 0 0 0 +SCALE 3 .3 3 + +// Floor +OBJECT 1 +cube +material 1 +TRANS 0 -0.5 0 +ROTAT 0 0 0 +SCALE 10 1 10 + +// Ceiling +OBJECT 2 +cube +material 1 +TRANS 0 10.5 0 +ROTAT 0 0 90 +SCALE 1 10 10 + +// Back wall +OBJECT 3 +cube +material 1 +TRANS 0 5 -5.5 +ROTAT 0 90 0 +SCALE 1 10 10 + +// Left wall +OBJECT 4 +cube +material 2 +TRANS -5.5 5 0 +ROTAT 0 0 0 +SCALE 1 10 10 + +// Right wall +OBJECT 5 +cube +material 3 +TRANS 5.5 5 0 +ROTAT 0 0 0 +SCALE 1 10 10 + +// Sphere +OBJECT 6 +sphere +material 7 +TRANS -3 4 -1 +ROTAT 0 0 0 +SCALE 3 3 3 + +// Back light +OBJECT 7 +cube +material 4 +TRANS 0 5 -5 +ROTAT 0 0 0 +SCALE 1 7 .3 + +// Sphere +OBJECT 8 +sphere +material 5 +TRANS 3 6 -2 +ROTAT 0 0 0 +SCALE 3 3 3 + +// Cube +OBJECT 9 +cube +material 8 +TRANS -2 1 2 +ROTAT 45 45 0 +SCALE 1 1 1 + +// Sphere +OBJECT 10 +sphere +material 7 +TRANS 2 3 5 +ROTAT 0 0 0 +SCALE 1 1 1 + +// Sphere +OBJECT 11 +sphere +material 8 +TRANS 0.5 5 5 +ROTAT 0 0 0 +SCALE 2 2 2 + +// Sphere +OBJECT 12 +sphere +material 8 +TRANS 0 1 0 +ROTAT 0 0 0 +SCALE 1 1 1 \ No newline at end of file diff --git a/scenes/cornell2.txt b/scenes/cornell2.txt new file mode 100644 index 0000000..a7bc43b --- /dev/null +++ b/scenes/cornell2.txt @@ -0,0 +1,246 @@ +// Emissive material (light) +MATERIAL 0 +RGB 0 0 0 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// 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 + +// Emissive material (light) +MATERIAL 4 +RGB 1 1 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 30 + +// SSS +MATERIAL 5 +RGB .851 .569 .392 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 +SSS 1 + +// SSS Compare +MATERIAL 6 +RGB .851 .569 .392 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Diffuse black +MATERIAL 7 +RGB .05 .05 .05 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Camera +CAMERA +RES 800 800 +FOVY 45 +ITERATIONS 5000 +DEPTH 8 +FILE cornell +EYE 0.0 5 14 +VIEW 0 0 -1 +UP 0 1 0 +DOF 9 + + +// Ceiling light +OBJECT 0 +cube +material 0 +TRANS 0 10 5 +ROTAT 0 0 0 +SCALE 3 .3 3 + +// Floor +OBJECT 1 +cube +material 7 +TRANS 0 -0.5 0 +ROTAT 0 0 0 +SCALE 10 1 10 + +// Ceiling +OBJECT 2 +cube +material 7 +TRANS 0 10.5 0 +ROTAT 0 0 90 +SCALE 1 10 10 + +// Back wall +OBJECT 3 +cube +material 7 +TRANS 0 5 -5.5 +ROTAT 0 90 0 +SCALE 1 10 10 + +// Left wall +OBJECT 4 +cube +material 7 +TRANS -5.5 5 0 +ROTAT 0 0 0 +SCALE 1 10 10 + +// Right wall +OBJECT 5 +cube +material 7 +TRANS 5.5 5 0 +ROTAT 0 0 0 +SCALE 1 10 10 + +// Back light top +OBJECT 6 +cube +material 4 +TRANS 0 7 -5 +ROTAT 0 0 0 +SCALE 2 2 .3 + +// Sphere SSS +OBJECT 7 +sphere +material 5 +TRANS 2 5 3 +ROTAT 0 0 0 +SCALE 2 2 2 + +// Sphere +OBJECT 8 +sphere +material 6 +TRANS -2 5 3 +ROTAT 0 0 0 +SCALE 2 2 2 + +// Sphere SSS +OBJECT 9 +sphere +material 5 +TRANS 0 3 -3 +ROTAT 0 0 0 +SCALE 2 2 2 + +// Sphere +OBJECT 10 +sphere +material 6 +TRANS 0 7 -3 +ROTAT 0 0 0 +SCALE 2 2 2 + +// Back light bottom +OBJECT 11 +cube +material 4 +TRANS 0 3 -5 +ROTAT 0 0 0 +SCALE 2 2 .3 + +// Floor light +OBJECT 12 +cube +material 0 +TRANS 0 0 5 +ROTAT 0 0 0 +SCALE 3 .3 3 + + +// Back light left +OBJECT 13 +cube +material 4 +TRANS -2 5 -5 +ROTAT 0 0 0 +SCALE 2 2 .3 + +// Back light right +OBJECT 14 +cube +material 4 +TRANS 2 5 -5 +ROTAT 0 0 0 +SCALE 2 2 .3 + +// Sphere SSS +OBJECT 15 +sphere +material 5 +TRANS 2 3 0 +ROTAT 0 0 0 +SCALE 2 2 2 + +// Sphere +OBJECT 16 +sphere +material 6 +TRANS -2 7 0 +ROTAT 0 0 0 +SCALE 2 2 2 + +// Back light diag left +OBJECT 17 +cube +material 4 +TRANS -2 7 -5 +ROTAT 0 0 0 +SCALE 2 2 .3 + +// Back light diag right +OBJECT 18 +cube +material 4 +TRANS 2 3 -5 +ROTAT 0 0 0 +SCALE 2 2 .3 \ No newline at end of file diff --git a/scenes/cornell3.txt b/scenes/cornell3.txt new file mode 100644 index 0000000..f2a0dd3 --- /dev/null +++ b/scenes/cornell3.txt @@ -0,0 +1,192 @@ +// Emissive material (light) +MATERIAL 0 +RGB 1 1 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 3 + +// 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 + +// Emissive material (light) +MATERIAL 4 +RGB 1 1 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 40 + +// SSS +MATERIAL 5 +RGB .902 .894 .847 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 +SSS 1 + +// SSS with reflection +MATERIAL 6 +RGB .149 .38 .612 +SPECEX 1000 +SPECRGB .149 .38 .612 +REFL 1 +REFR 0 +REFRIOR 1.486 +EMITTANCE 0 +SSS 1 + +// Pure reflection +MATERIAL 7 +RGB .149 .38 .612 +SPECEX 1000 +SPECRGB .149 .38 .612 +REFL 1 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Diffuse black +MATERIAL 8 +RGB .3 .3 .3 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Camera +CAMERA +RES 800 800 +FOVY 45 +ITERATIONS 5000 +DEPTH 8 +FILE cornell +EYE 0.0 5 14 +VIEW 0 0 -1 +UP 0 1 0 +DOF 9 + + +// Ceiling light +OBJECT 0 +cube +material 0 +TRANS 0 10 5 +ROTAT 0 0 0 +SCALE 3 .3 3 + +// Floor +OBJECT 1 +cube +material 8 +TRANS 0 -0.5 0 +ROTAT 0 0 0 +SCALE 10 1 10 + +// Ceiling +OBJECT 2 +cube +material 8 +TRANS 0 10.5 0 +ROTAT 0 0 90 +SCALE 1 10 10 + +// Back wall +OBJECT 3 +cube +material 8 +TRANS 0 5 -5.5 +ROTAT 0 90 0 +SCALE 1 10 10 + +// Left wall +OBJECT 4 +cube +material 2 +TRANS -5.5 5 0 +ROTAT 0 0 0 +SCALE 1 10 10 + +// Right wall +OBJECT 5 +cube +material 3 +TRANS 5.5 5 0 +ROTAT 0 0 0 +SCALE 1 10 10 + +// Sphere SSS w/ reflection +OBJECT 6 +sphere +material 6 +TRANS -1.5 5 3 +ROTAT 0 0 0 +SCALE 2 2 2 + +// Sphere pure reflection +OBJECT 7 +sphere +material 7 +TRANS 1.5 5 3 +ROTAT 0 0 0 +SCALE 2 2 2 + +// Back light left +OBJECT 8 +sphere +material 4 +TRANS -1.5 5 2 +ROTAT 0 0 0 +SCALE 1.5 1.5 .1 + +// Back light right +OBJECT 9 +sphere +material 4 +TRANS 1.5 5 2 +ROTAT 0 0 0 +SCALE 1.5 1.5 .1 + +// Blocking cube +OBJECT 10 +cube +material 8 +TRANS 0 5 0 +ROTAT 0 0 0 +SCALE 0.4 8 10 \ No newline at end of file diff --git a/scenes/cornell4.txt b/scenes/cornell4.txt new file mode 100644 index 0000000..bf9d0fa --- /dev/null +++ b/scenes/cornell4.txt @@ -0,0 +1,246 @@ +// Emissive material (light) +MATERIAL 0 +RGB 0 0 0 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// 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 + +// Emissive material (light) +MATERIAL 4 +RGB 1 1 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 30 + +// SSS +MATERIAL 5 +RGB .851 .569 .392 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 +SSS 1 + +// SSS Compare +MATERIAL 6 +RGB .851 .569 .392 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Diffuse black +MATERIAL 7 +RGB .05 .05 .05 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Camera +CAMERA +RES 800 800 +FOVY 45 +ITERATIONS 5000 +DEPTH 8 +FILE cornell +EYE 0.0 5 14 +VIEW 0 0 -1 +UP 0 1 0 +DOF 9 + + +// Ceiling light +OBJECT 0 +cube +material 0 +TRANS 0 10 5 +ROTAT 0 0 0 +SCALE 3 .3 3 + +// Floor +OBJECT 1 +cube +material 7 +TRANS 0 -0.5 0 +ROTAT 0 0 0 +SCALE 10 1 10 + +// Ceiling +OBJECT 2 +cube +material 7 +TRANS 0 10.5 0 +ROTAT 0 0 90 +SCALE 1 10 10 + +// Back wall +OBJECT 3 +cube +material 7 +TRANS 0 5 -5.5 +ROTAT 0 90 0 +SCALE 1 10 10 + +// Left wall +OBJECT 4 +cube +material 7 +TRANS -5.5 5 0 +ROTAT 0 0 0 +SCALE 1 10 10 + +// Right wall +OBJECT 5 +cube +material 7 +TRANS 5.5 5 0 +ROTAT 0 0 0 +SCALE 1 10 10 + +// Back light top +OBJECT 6 +cube +material 4 +TRANS 0 7 -5 +ROTAT 0 0 0 +SCALE 2 2 .3 + +// Sphere SSS +OBJECT 7 +cube +material 5 +TRANS 2 5 3 +ROTAT 45 45 0 +SCALE 2 2 2 + +// Sphere +OBJECT 8 +cube +material 6 +TRANS -2 5 3 +ROTAT 45 45 0 +SCALE 2 2 2 + +// Sphere SSS +OBJECT 9 +cube +material 5 +TRANS 0 3 -3 +ROTAT 45 45 0 +SCALE 2 2 2 + +// Sphere +OBJECT 10 +cube +material 6 +TRANS 0 7 -3 +ROTAT 45 45 0 +SCALE 2 2 2 + +// Back light bottom +OBJECT 11 +cube +material 4 +TRANS 0 3 -5 +ROTAT 0 0 0 +SCALE 2 2 .3 + +// Floor light +OBJECT 12 +cube +material 0 +TRANS 0 0 5 +ROTAT 0 0 0 +SCALE 3 .3 3 + + +// Back light left +OBJECT 13 +cube +material 4 +TRANS -2 5 -5 +ROTAT 0 0 0 +SCALE 2 2 .3 + +// Back light right +OBJECT 14 +cube +material 4 +TRANS 2 5 -5 +ROTAT 0 0 0 +SCALE 2 2 .3 + +// Sphere SSS +OBJECT 15 +cube +material 5 +TRANS 2 3 0 +ROTAT 45 45 0 +SCALE 2 2 2 + +// Sphere +OBJECT 16 +cube +material 6 +TRANS -2 7 0 +ROTAT 45 45 0 +SCALE 2 2 2 + +// Back light diag left +OBJECT 17 +cube +material 4 +TRANS -2 7 -5 +ROTAT 0 0 0 +SCALE 2 2 .3 + +// Back light diag right +OBJECT 18 +cube +material 4 +TRANS 2 3 -5 +ROTAT 0 0 0 +SCALE 2 2 .3 \ No newline at end of file diff --git a/scenes/cornell5.txt b/scenes/cornell5.txt new file mode 100644 index 0000000..4d2ca60 --- /dev/null +++ b/scenes/cornell5.txt @@ -0,0 +1,214 @@ +// Emissive material (light) +MATERIAL 0 +RGB 1 1 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 2 + +// Diffuse white +MATERIAL 1 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Diffuse red +MATERIAL 2 +RGB .85 .35 .35 +SPECEX 30 +SPECRGB 1 1 1 +REFL 1 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Diffuse green +MATERIAL 3 +RGB .35 .85 .35 +SPECEX 1000 +SPECRGB 1 1 1 +REFL 1 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Emissive material (light) +MATERIAL 4 +RGB 1 1 0 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 2 + +// Glossy material +MATERIAL 5 +RGB .35 .35 .85 +SPECEX 5 +SPECRGB .8 .8 .8 +REFL 1 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Glossy material +MATERIAL 6 +RGB .35 .85 .35 +SPECEX 0 +SPECRGB .8 .8 .8 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Mirror white +MATERIAL 7 +RGB .98 .98 .98 +SPECEX 100 +SPECRGB 1 1 1 +REFL 1 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Refraction +MATERIAL 8 +RGB .847 .949 .9529 +SPECEX 1000 +SPECRGB 1 1 1 +REFL 0 +REFR 1 +REFRIOR 1.309 +EMITTANCE 0 + +// Camera +CAMERA +RES 800 800 +FOVY 45 +ITERATIONS 5000 +DEPTH 8 +FILE cornell +EYE 0.0 5 14 +VIEW 0 0 -1 +UP 0 1 0 +DOF 9 + + +// Ceiling light +OBJECT 0 +cube +material 0 +TRANS 0 10 0 +ROTAT 0 0 0 +SCALE 3 .3 3 + +// Floor +OBJECT 1 +cube +material 1 +TRANS 0 -0.5 0 +ROTAT 0 0 0 +SCALE 10 1 28 + +// Ceiling +OBJECT 2 +cube +material 1 +TRANS 0 10.5 0 +ROTAT 0 0 90 +SCALE 1 10 28 + +// Back wall +OBJECT 3 +cube +material 1 +TRANS 0 5 -5.5 +ROTAT 0 90 0 +SCALE 1 10 10 + +// Left wall +OBJECT 4 +cube +material 2 +TRANS -5.5 5 0 +ROTAT 0 0 0 +SCALE 1 10 28 + +// Right wall +OBJECT 5 +cube +material 3 +TRANS 5.5 5 0 +ROTAT 0 0 0 +SCALE 1 10 28 + +// Sphere +OBJECT 6 +sphere +material 7 +TRANS -3 4 -1 +ROTAT 0 0 0 +SCALE 3 3 3 + +// Back light +OBJECT 7 +cube +material 4 +TRANS 0 5 -5 +ROTAT 0 0 0 +SCALE 1 7 .3 + +// Sphere +OBJECT 8 +sphere +material 5 +TRANS 3 6 -2 +ROTAT 0 0 0 +SCALE 3 3 3 + +// Cube +OBJECT 9 +cube +material 8 +TRANS -2 1 2 +ROTAT 45 45 0 +SCALE 1 1 1 + +// Sphere +OBJECT 10 +sphere +material 7 +TRANS 2 3 5 +ROTAT 0 0 0 +SCALE 1 1 1 + +// Sphere +OBJECT 11 +sphere +material 8 +TRANS 0.5 5 5 +ROTAT 0 0 0 +SCALE 2 2 2 + +// Sphere +OBJECT 12 +sphere +material 8 +TRANS 0 1 0 +ROTAT 0 0 0 +SCALE 1 1 1 + +// Front wall +OBJECT 13 +cube +material 1 +TRANS 0 5 14.5 +ROTAT 0 90 0 +SCALE 1 10 10 \ No newline at end of file diff --git a/scenes/cornell6.txt b/scenes/cornell6.txt new file mode 100644 index 0000000..a062302 --- /dev/null +++ b/scenes/cornell6.txt @@ -0,0 +1,136 @@ +// 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 30 +SPECRGB 1 1 1 +REFL 1 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Diffuse green +MATERIAL 3 +RGB .35 .85 .35 +SPECEX 0 +SPECRGB 1 1 1 +REFL 1 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Emissive material (light) +MATERIAL 4 +RGB 1 1 0 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 1 + +// Test material +MATERIAL 5 +RGB .35 .35 .85 +SPECEX 5 +SPECRGB .8 .8 .8 +REFL 1 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Camera +CAMERA +RES 200 200 +FOVY 45 +ITERATIONS 5000 +DEPTH 8 +FILE cornell +EYE 0.0 5 14 +VIEW 0 0 -1 +UP 0 1 0 +DOF 9 + + +// Ceiling light +OBJECT 0 +cube +material 0 +TRANS 0 10 0 +ROTAT 0 0 0 +SCALE 3 .3 3 + +// Floor +OBJECT 1 +cube +material 1 +TRANS 0 -0.5 0 +ROTAT 0 0 0 +SCALE 10 1 10 + +// Ceiling +OBJECT 2 +cube +material 1 +TRANS 0 10.5 0 +ROTAT 0 0 90 +SCALE 1 10 10 + +// Back wall +OBJECT 3 +cube +material 1 +TRANS 0 5 -5.5 +ROTAT 0 90 0 +SCALE 1 10 10 + +// Left wall +OBJECT 4 +cube +material 2 +TRANS -5.5 5 0 +ROTAT 0 0 0 +SCALE 1 10 10 + +// Right wall +OBJECT 5 +cube +material 3 +TRANS 5.5 5 0 +ROTAT 0 0 0 +SCALE 1 10 10 + +// Sphere +OBJECT 6 +sphere +material 5 +TRANS 0 5 0 +ROTAT 0 0 0 +SCALE 4 4 4 + +// Back light +OBJECT 7 +cube +material 4 +TRANS 0 5 -5 +ROTAT 0 0 0 +SCALE 1 7 .3 \ No newline at end of file diff --git a/src/interactions.h b/src/interactions.h index f0bf0b6..85249f7 100644 --- a/src/interactions.h +++ b/src/interactions.h @@ -41,6 +41,103 @@ glm::vec3 calculateRandomDirectionInHemisphere( + sin(around) * over * perpendicularDirection2; } +/** +* Computes a cosine-weighted random direction in a hemisphere. +* Used for non-perfect specular lighting. +*/ +__host__ __device__ +glm::vec3 calculateRandomSpecDirection( +glm::vec3 normal, float n, thrust::default_random_engine &rng) { + thrust::uniform_real_distribution u01(0, 1); + + // Weighted cosine by specular exponent + float up = pow(u01(rng), 1/(n+1)); // cos(theta) + float over = sqrt(1 - up * up); // sin(theta) + float around = u01(rng) * TWO_PI; // phi + + // Find a direction that is not the normal based off of whether or not the + // normal's components are all equal to sqrt(1/3) or whether or not at + // least one component is less than sqrt(1/3). Learned this trick from + // Peter Kutz. + + glm::vec3 directionNotNormal; + if (abs(normal.x) < SQRT_OF_ONE_THIRD) { + directionNotNormal = glm::vec3(1, 0, 0); + } + else if (abs(normal.y) < SQRT_OF_ONE_THIRD) { + directionNotNormal = glm::vec3(0, 1, 0); + } + else { + directionNotNormal = glm::vec3(0, 0, 1); + } + + // Use not-normal direction to generate two perpendicular directions + glm::vec3 perpendicularDirection1 = + 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; +} + +/** +* Computes a cosine-weighted random direction and a guessed new origin in a hemisphere for SSS +* Used for SSS. +* Really a simplified version of Dipole approximation +*/ +__host__ __device__ +void calculateSSSOut(Ray &r, glm::vec3 &intersect, glm::vec3 &inDirection, glm::vec3 &normal, thrust::default_random_engine &rng) { + // Punch in a "pinhole" along negative normal direction + // Use endpoint as origin, and shoot a random ray toward surface from inside + glm::vec3 origin = intersect + normalize(inDirection)*0.0001f - normal; + glm::vec3 scatterDirection = calculateRandomDirectionInHemisphere(normal, rng); + // Shoot the ray from inside + // Approximate scattered origin with normal length + glm::vec3 scatterOrigin = origin + normalize(scatterDirection) * glm::length(normal); + // Does not necessarily put the ray outside geometry; if not it will be catched at the next intersection test + // Passing geometry all the way to this function is REALLY slow; therefore the hacky way + r.origin = scatterOrigin + normalize(scatterDirection)*0.0005f; + r.direction = scatterDirection; +} + +/** +* Computes refraction direction +* Used for refraction lighting. +*/ +__host__ __device__ +glm::vec3 calculateRefractDirection(bool outside, glm::vec3 inDirection, glm::vec3 normal, float n1, float n2) { + float mu; + glm::vec3 n = normalize(normal); + glm::vec3 i = normalize(inDirection); + if (outside){ + // Enter medium, leave air + mu = n1/n2; + } + else { + // Enter air, leave medium + mu = n2/n1; + } + float cosI = -dot(n, i); + float cosTsq = 1 - mu*mu*(1 - cosI*cosI); + if (cosTsq < 0){ + return glm::vec3(0.0f); + } + else { + return (mu * i) + (mu * cosI - sqrt(cosTsq)) * n; + } +} + +/** +* Computes the perfect reflection direction on a given normal +* Used for perfect specular lighting. (not used for non-perfect specular lighting) +*/ +__host__ __device__ +glm::vec3 calculatePerfectSpecDirection(glm::vec3 inDirection, glm::vec3 normal) { + return inDirection - 2.0f * glm::dot(inDirection, normal) * normal; +} + /** * Scatter a ray with some probabilities according to the material properties. * For example, a diffuse surface scatters in a cosine-weighted hemisphere. @@ -61,15 +158,127 @@ glm::vec3 calculateRandomDirectionInHemisphere( * * 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, + PathRay &path, 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 range(0, 1); + // Diffuse is for sure to happen + glm::vec3 diffuseDirection = calculateRandomDirectionInHemisphere(path.normal, rng); + glm::vec3 diffuseColor = path.color * m.color; + + if (m.hasRefractive == 1.0f){ + float n1 = 1; + float n2 = m.indexOfRefraction; + float r0 = pow((n1-n2)/(n1+n2), 2); + + //glm::vec3 perfectSpecDirection = calculatePerfectSpecDirection(inDirection, normal); + + // Jitter intersection normal to get jittered specular reflections + glm::vec3 glossNormal = calculateRandomSpecDirection(path.normal, m.specular.exponent, rng); + glm::vec3 glossDirection = calculatePerfectSpecDirection(path.ray.direction, glossNormal); + glm::vec3 glossColor = path.color * m.specular.color; + + glm::vec3 refractDirection = calculateRefractDirection(path.outside, path.ray.direction, path.normal, n1, n2); + + // Schlick coefficient + float r = r0 + (1 - r0)*pow((1 - dot(normalize(glossDirection), normalize(path.normal))), 5); + + if (refractDirection == glm::vec3(0.0f)){ + // Total internal reflection + path.color = glossColor * r; + path.ray.direction = glossDirection; + path.ray.origin = path.intersect; + } + else { + glossColor = glossColor * r; + float dI = glm::length(diffuseColor); + float split = dI / (dI + glm::length(glossColor)); + float pick = range(rng); + if (pick < split){ + path.color = diffuseColor / split; + // Move on the original incident direction + path.ray.origin = path.intersect + normalize(path.ray.direction)*0.0005f; + // Update to outbound direction + path.ray.direction = refractDirection; + } + else { + path.color = glossColor / split; + path.ray.direction = glossDirection; + path.ray.origin = path.intersect; + } + } + } + else if (m.hasReflective == 1.0f){ + //glm::vec3 perfectSpecDirection = calculatePerfectSpecDirection(inDirection, normal); + glm::vec3 glossNormal = calculateRandomSpecDirection(path.normal, m.specular.exponent, rng); + glm::vec3 glossDirection = calculatePerfectSpecDirection(path.ray.direction, glossNormal); + glm::vec3 glossColor = path.color * m.specular.color; + + float dI = glm::length(diffuseColor); + float split = dI / (dI + glm::length(glossColor)); + float pick = range(rng); + + if (pick < split){ + /* + color = diffuseColor / split; + ray.direction = diffuseDirection; + */ + if (m.hasSSS == 1.0f){ + if (path.outside){ + calculateSSSOut(path.ray, path.intersect, path.ray.direction, path.normal, rng); + } + else { + float pick = range(rng); + if (pick < 0.5){ + path.ray.origin = path.intersect; + path.ray.direction = diffuseDirection; + } + else { + path.ray.origin = path.intersect + normalize(path.ray.direction)*0.0005f; + path.ray.direction = path.ray.direction; + } + path.color = diffuseColor / split / 2.0f; + } + path.color = diffuseColor / split; + } + else { + path.color = diffuseColor / split; + path.ray.origin = path.intersect; + path.ray.direction = diffuseDirection; + } + } + else { + path.color = glossColor / split; + path.ray.origin = path.intersect; + path.ray.direction = glossDirection; + } + } + else if (m.hasSSS == 1.0f){ + if (path.outside){ + // Only scatter under surface if it's an incoming light + calculateSSSOut(path.ray, path.intersect, path.ray.direction, path.normal, rng); + } + else { + // Otherwise light is going out; should be the offset diffuse reflection + float pick = range(rng); + if (pick < 0.5){ + path.ray.origin = path.intersect; + path.ray.direction = diffuseDirection; + } + else { + path.ray.origin = path.intersect + normalize(path.ray.direction)*0.0005f; + path.ray.direction = path.ray.direction; + } + } + path.color = diffuseColor * 2.0f; + } + else { + path.color = diffuseColor; + path.ray.origin = path.intersect; + path.ray.direction = diffuseDirection; + } } diff --git a/src/intersections.h b/src/intersections.h index 5f3613d..ce01f57 100644 --- a/src/intersections.h +++ b/src/intersections.h @@ -36,109 +36,108 @@ __host__ __device__ glm::vec3 multiplyMV(glm::mat4 m, glm::vec4 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. - */ -__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; +* 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. +*/ +__host__ __device__ float boxIntersectionTest(Geom &box, Ray &r, + glm::vec3 &intersectionPoint, glm::vec3 &normal, bool &outside) { + 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))); + + 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. - */ -__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); -} +* 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. +*/ +__host__ __device__ float sphereIntersectionTest(Geom &sphere, Ray &r, + glm::vec3 &intersectionPoint, glm::vec3 &normal, bool &outside) { + 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.0f)); + normal = glm::normalize(multiplyMV(sphere.invTranspose, glm::vec4(objspaceIntersection, 0.0f))); + 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..d347bcc 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -1,6 +1,8 @@ #include "main.h" #include "preview.h" #include +#include "glm/glm.hpp" +#include "glm/gtx/norm.hpp" static std::string startTimeString; static bool camchanged = false; @@ -80,6 +82,20 @@ void runCuda() { cam.view = glm::vec3(rotmat * glm::vec4(v, 0.f)); cam.up = glm::vec3(rotmat * glm::vec4(u, 0.f)); cam.position += cammove.x * r + cammove.y * u + cammove.z * v; + + // Camera to grid center + float distance = cam.resolution.x / 2 / tan(cam.fov.x / 2); + cam.toGrid = glm::vec3(cam.view.x*distance, cam.view.y*distance, cam.view.z*distance); + // Find camera right vector + float rAngle = -PI / 2; + float qx = cam.view.x * sin(rAngle / 2); + float qy = cam.view.y * sin(rAngle / 2); + float qz = cam.view.z * sin(rAngle / 2); + float qw = cos(rAngle / 2); + + glm::quat q = glm::quat(qw, qx, qy, qz); + cam.right = q * cam.up; + theta = phi = 0; cammove = glm::vec3(); camchanged = false; diff --git a/src/pathtrace.cu b/src/pathtrace.cu index 7fbaf30..4b9a6aa 100644 --- a/src/pathtrace.cu +++ b/src/pathtrace.cu @@ -4,6 +4,11 @@ #include #include #include +#include +#include +#include + +#include #include "sceneStructs.h" #include "scene.h" @@ -14,9 +19,16 @@ #include "intersections.h" #include "interactions.h" +#define ANTIALIASING 0 +#define DOF 0 +#define USETHRUSTCOMPACTION 0 +#define TIMEFUNCTION 0 #define FILENAME (strrchr(__FILE__, '/') ? strrchr(__FILE__, '/') + 1 : __FILE__) #define checkCUDAError(msg) checkCUDAErrorFn(msg, FILENAME, __LINE__) +#define ERRORCHECK 0 void checkCUDAErrorFn(const char *msg, const char *file, int line) { +#if ERRORCHECK + + cudaDeviceSynchronize(); cudaError_t err = cudaGetLastError(); if (cudaSuccess == err) { return; @@ -27,7 +39,8 @@ void checkCUDAErrorFn(const char *msg, const char *file, int line) { fprintf(stderr, " (%s:%d)", file, line); } fprintf(stderr, ": %s: %s\n", msg, cudaGetErrorString(err)); - exit(EXIT_FAILURE); + exit(EXIT_FAILURE); +#endif ERRORCHECK } __host__ __device__ thrust::default_random_engine random_engine( @@ -59,9 +72,17 @@ __global__ void sendImageToPBO(uchar4* pbo, glm::ivec2 resolution, } static Scene *hst_scene = NULL; -static glm::vec3 *dev_image = NULL; -// TODO: static variables for device memory, scene/camera info, etc -// ... +__constant__ static glm::vec3 *dev_image = NULL; +__constant__ static Geom* dev_geoms = NULL; +__constant__ static Material* dev_materials = NULL; +__constant__ static glm::vec3 *dev_oversample_image = NULL; +static int geomcount = 0; +static int oversampling_pass = 3; +// Temp variables for stream compaction +__constant__ static int *dv_f_tmp = NULL; +__constant__ static int *dv_idx_tmp = NULL; +__constant__ static PathRay *dv_out_tmp = NULL; +__constant__ static int *dv_c_tmp = NULL; void pathtraceInit(Scene *scene) { hst_scene = scene; @@ -70,42 +91,235 @@ void pathtraceInit(Scene *scene) { 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_oversample_image, pixelcount * sizeof(glm::vec3)); + + Geom* hst_geoms = hst_scene->geoms.data(); + Material* hst_materials = hst_scene->materials.data(); + + geomcount = hst_scene->geoms.size(); + + cudaMalloc((void**)&dev_geoms, hst_scene->geoms.size()*sizeof(Geom)); + cudaMalloc((void**)&dev_materials, hst_scene->materials.size()*sizeof(Material)); + cudaMemcpy(dev_geoms, hst_geoms, hst_scene->geoms.size()*sizeof(Geom), cudaMemcpyHostToDevice); + cudaMemcpy(dev_materials, hst_materials, hst_scene->materials.size()*sizeof(Material), cudaMemcpyHostToDevice); + + // Allocate temp vars + cudaMalloc((void**)&dv_f_tmp, pixelcount*sizeof(int)); + cudaMalloc((void**)&dv_idx_tmp, pixelcount*sizeof(int)); + cudaMalloc((void**)&dv_out_tmp, pixelcount*sizeof(PathRay)); + cudaMalloc((void**)&dv_c_tmp, sizeof(int)); checkCUDAError("pathtraceInit"); } void pathtraceFree() { cudaFree(dev_image); // no-op if dev_image is null - // TODO: clean up the above static variables + cudaFree(dev_geoms); + cudaFree(dev_materials); + cudaFree(dev_oversample_image); + cudaFree(dv_f_tmp); + cudaFree(dv_idx_tmp); + cudaFree(dv_out_tmp); + cudaFree(dv_c_tmp); checkCUDAError("pathtraceFree"); } -/** - * Example function to generate static and test the CUDA-GL interop. - * Delete this once you're done looking at it! - */ -__global__ void generateNoiseDeleteMe(Camera cam, int iter, glm::vec3 *image) { - int x = (blockIdx.x * blockDim.x) + threadIdx.x; - int y = (blockIdx.y * blockDim.y) + threadIdx.y; +__global__ void initRayGrid(PathRay *oGrid, const Camera cam){ + // From camera as single point, to image grid with FOV + 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); + if (x < cam.resolution.x && y < cam.resolution.y) { + int index = x + (y * cam.resolution.x); + //PathRay pr; + oGrid[index].index = index; + oGrid[index].color = glm::vec3(1.0f); - thrust::default_random_engine rng = random_engine(iter, index, 0); - thrust::uniform_real_distribution u01(0, 1); + oGrid[index].ray.origin = cam.position; + oGrid[index].terminate = false; + oGrid[index].matId = -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)); - } + // Grid center to pixel + float pX = x - cam.resolution.x / 2; + float pY = cam.resolution.y / 2 - y; + + // Vector: grid center to pixel + glm::vec3 o2px = cam.right*pX + cam.up*pY; + //glm::vec3 o2px = glm::vec3(cam.right.x*pX + cam.up.x*pY, cam.right.y*pX + cam.up.y*pY, cam.right.z*pX + cam.up.z*pY); + // Ray vector + oGrid[index].ray.direction = cam.toGrid + o2px; + + //oGrid[index] = pr; + + // Ray direction debug + //float l = glm::length(ray.ray.direction); + //image[index] += glm::vec3(abs(ray.ray.direction.x / l), abs(ray.ray.direction.y / l), 0); + } +} + + +__global__ void interesect(PathRay * grid, const Geom * iGeoms, const Camera cam, const int grid_size, const int geomcount){ + // From camera as single point, to image grid with FOV + extern __shared__ Geom geoms[]; + + int bIndex = threadIdx.x + threadIdx.y * blockDim.x; + int x = (blockIdx.x * blockDim.x) + threadIdx.x; + int y = (blockIdx.y * blockDim.y) + threadIdx.y; + int index = x + (y * cam.resolution.x); + if (bIndex < geomcount){ + geoms[bIndex] = iGeoms[bIndex]; + } + + __syncthreads(); + + if (index < grid_size) { + // Intersection test + //PathRay pr = grid[index]; + grid[index].hasIntersect = false; + glm::vec3 iPoint; + glm::vec3 iNormal; + bool outside = false; + + float rayLength = 0; + float oldLength = -1; + + for (int i = 0; i < geomcount; ++i){ + if (geoms[i].type == SPHERE){ + rayLength = sphereIntersectionTest(geoms[i], grid[index].ray, iPoint, iNormal, outside); + } + else { + rayLength = boxIntersectionTest(geoms[i], grid[index].ray, iPoint, iNormal, outside); + } + // Find the nearest intersection + if (rayLength > -1){ + grid[index].hasIntersect = true; + if (oldLength == -1 || rayLength < oldLength){ + grid[index].intersect = iPoint; + grid[index].normal = iNormal; + grid[index].outside = outside; + grid[index].matId = geoms[i].materialid; + oldLength = rayLength; + } + } + /* + Geom g = iGeoms[i]; + if (g.type == SPHERE){ + rayLength = sphereIntersectionTest(g, pr.ray, iPoint, iNormal, outside); + } else { + rayLength = boxIntersectionTest(g, pr.ray, iPoint, iNormal, outside); + } + // Find the nearest intersection + if (rayLength != -1.0f){ + pr.hasIntersect = true; + if (oldLength == -1.0f || rayLength < oldLength){ + oldLength = rayLength; + pr.intersect = iPoint; + pr.normal = iNormal; + pr.matId = g.materialid; + pr.outside = outside; + } + } + */ + } + //grid[index] = pr; + } +}; + +__global__ void scatter(PathRay *grid, const Material *iMaterials, const Camera cam, const int grid_size, const int iter, const int depth){ + // From camera as single point, to image grid with FOV + int x = (blockIdx.x * blockDim.x) + threadIdx.x; + int y = (blockIdx.y * blockDim.y) + threadIdx.y; + + int index = x + (y * cam.resolution.x); + + __syncthreads(); + + if (index < grid_size) { + //PathRay pr = grid[index]; + //Material m = iMaterials[pr.matId]; + scatterRay(grid[index], iMaterials[grid[index].matId], random_engine(iter, index, depth)); + //grid[index] = pr; + } +}; + +__global__ void terminatePath(PathRay *grid, const Material *iMaterials, const Camera cam, const int grid_size){ + // From camera as single point, to image grid with FOV + int x = (blockIdx.x * blockDim.x) + threadIdx.x; + int y = (blockIdx.y * blockDim.y) + threadIdx.y; + + int index = x + (y * cam.resolution.x); + + if (index < grid_size) { + if (grid[index].hasIntersect){ + // Hits a light + if (iMaterials[grid[index].matId].emittance > 0.0f){ + grid[index].terminate = true; + grid[index].color = grid[index].color * iMaterials[grid[index].matId].color * iMaterials[grid[index].matId].emittance; + } + } + else { + // No intersections + grid[index].terminate = true; + grid[index].color = glm::vec3(0.0f); + } + //grid[index] = pr; + } +} + +__global__ void fillPixel(const PathRay *grid, glm::vec3 *image, const Camera cam, const int grid_size){ + // From camera as single point, to image grid with FOV + int x = (blockIdx.x * blockDim.x) + threadIdx.x; + int y = (blockIdx.y * blockDim.y) + threadIdx.y; + + int index = x + (y * cam.resolution.x); + + if (index < grid_size) { + PathRay pr = grid[index]; + if (pr.terminate){ + image[pr.index] += pr.color; + } + } +} + +#if ANTIALIASING +__global__ void avgOversample(glm::vec3 *oImage, const glm::vec3 *tempImage, const Camera cam, const int passes){ + // From camera as single point, to image grid with FOV + 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); + oImage[index] += tempImage[index] / (float)passes; + } +} + +__global__ void jitterRay(PathRay *grid, thrust::default_random_engine rng, const Camera cam){ + // From camera as single point, to image grid with FOV + 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); + PathRay pr = grid[index]; +#if DOF + thrust::uniform_real_distribution u01(-0.1, 0.1); + // Find intersection to focal plane + glm::vec3 p = pr.ray.origin + glm::normalize(pr.ray.direction) * cam.dof; + glm::vec3 jitter = cam.up*u01(rng) + cam.right*u01(rng); + // Jitter ray origin + pr.ray.origin = pr.ray.origin + jitter; + // Update ray direction based on new origin + // So new ray always points to exact pixel cell on focal plane + pr.ray.direction = p - pr.ray.origin; +#else + thrust::uniform_real_distribution u01(-0.01, 0.01); + pr.ray.origin = glm::vec3(pr.ray.origin.x + u01(rng), pr.ray.origin.y + u01(rng), pr.ray.origin.z + u01(rng)); +#endif DOF + grid[index] = pr; + } } +#endif ANTIALIASING /** * Wrapper for the __global__ call that sets up the kernel calls and does a ton @@ -116,9 +330,9 @@ void pathtrace(uchar4 *pbo, int frame, int iter) { const Camera &cam = hst_scene->state.camera; const int pixelcount = cam.resolution.x * cam.resolution.y; - const int blockSideLength = 8; + const int blockSideLength = 8; const dim3 blockSize(blockSideLength, blockSideLength); - const dim3 blocksPerGrid( + dim3 blocksPerGrid( (cam.resolution.x + blockSize.x - 1) / blockSize.x, (cam.resolution.y + blockSize.y - 1) / blockSize.y); @@ -147,9 +361,87 @@ void pathtrace(uchar4 *pbo, int frame, int iter) { // * Finally, handle all of the paths that still haven't terminated. // (Easy way is to make them black or background-colored.) - // TODO: perform one iteration of path tracing + // Perform one iteration of path tracing + + thrust::device_vector dev_grid(pixelcount); + PathRay *dev_grid_ptr = thrust::raw_pointer_cast(&dev_grid[0]); + +#if ANTIALIASING + cudaMemset(dev_oversample_image, 0, pixelcount * sizeof(glm::vec3)); + for (int a = 0; a < oversampling_pass; a++){ + dev_grid.resize(pixelcount); + // initRayGrid + initRayGrid << > >(dev_grid_ptr, cam); + int grid_size = dev_grid.size(); + // Jitter for antialiasing oversampling; also accounts for DOF effect if enabled + jitterRay << > >(dev_grid_ptr, random_engine(iter, 0, oversampling_pass), cam); +#else + // initRayGrid + initRayGrid << > >(dev_grid_ptr, cam); + int grid_size = dev_grid.size(); +#endif ANTIALIASING + // For each traceDepth + for (int d = 0; d < traceDepth; d++){ + +#if TIMEFUNCTION + cudaEvent_t start, stop; + cudaEventCreate(&start); + cudaEventCreate(&stop); + cudaEventRecord(start); +#endif TIMEFUNCTION + + // Intersection test + interesect << > >(dev_grid_ptr, dev_geoms, cam, grid_size, geomcount); + checkCUDAError("intersect"); + + // Mark all terminated paths + terminatePath << > >(dev_grid_ptr, dev_materials, cam, grid_size); + checkCUDAError("terminatePath"); + + // Paint image +#if ANTIALIASING + fillPixel << > >(dev_grid_ptr, dev_oversample_image, cam, grid_size); +#else + fillPixel << > >(dev_grid_ptr, dev_image, cam, grid_size); +#endif ANTIALIASING + checkCUDAError("fillPixel"); + + // Stream compaction +#if USETHRUSTCOMPACTION + thrust::detail::normal_iterator> newGridEnd = thrust::remove_if(dev_grid.begin(), dev_grid.end(), is_terminated()); + checkCUDAError("thrustCompact"); + dev_grid.erase(newGridEnd, dev_grid.end()); + grid_size = dev_grid.size(); +#else + StreamCompaction::Efficient::compact(grid_size, dv_f_tmp, dv_idx_tmp, dv_out_tmp, dev_grid_ptr, dv_c_tmp); + cudaMemcpy(&grid_size, dv_c_tmp, sizeof(int), cudaMemcpyDeviceToHost); + cudaMemcpy(dev_grid_ptr, dv_out_tmp, grid_size * sizeof(PathRay), cudaMemcpyDeviceToDevice); + checkCUDAError("efficientCompact"); +#endif USETHRUSTCOMPACTION + + // Scatter + scatter << > >(dev_grid_ptr, dev_materials, cam, grid_size, iter, d); + checkCUDAError("scatter"); + + // Stream compaction stat + //printf("Iter: %d / Depth: %d / Grid size: %d\n", iter, d, grid_size); + +#if TIMEFUNCTION + cudaEventRecord(stop); + cudaEventSynchronize(stop); + float ms1 = 0; + cudaEventElapsedTime(&ms1, start, stop); - generateNoiseDeleteMe<<>>(cam, iter, dev_image); + // Exec time comparison: thrust vs. work-efficient + printf("Iter: %d / Depth: %d / Grid size: %d / Thrust?: %d / Time: %f\n", iter, d, grid_size, USETHRUSTCOMPACTION, ms1); +#endif TIMEFUNCTION + } + dev_grid.clear(); +#if ANTIALIASING + } + // Average oversampled colors and fill into image + avgOversample << > >(dev_image, dev_oversample_image, cam, oversampling_pass); +#endif ANTIALIASING /////////////////////////////////////////////////////////////////////////// diff --git a/src/scene.cpp b/src/scene.cpp index 5804ce3..f2b18b8 100644 --- a/src/scene.cpp +++ b/src/scene.cpp @@ -124,7 +124,10 @@ int Scene::loadCamera() { 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(), "DOF") == 0) { + camera.dof = atof(tokens[1].c_str()); + } utilityCore::safeGetline(fp_in, line); } @@ -135,6 +138,19 @@ int Scene::loadCamera() { float fovx = (atan(xscaled) * 180) / PI; camera.fov = glm::vec2(fovx, fovy); + // Camera to grid center + float distance = camera.resolution.x / 2 / tan(camera.fov.x / 2); + camera.toGrid = glm::vec3(camera.view.x*distance, camera.view.y*distance, camera.view.z*distance); + // Find camera right vector + float rAngle = -PI / 2; + float qx = camera.view.x * sin(rAngle / 2); + float qy = camera.view.y * sin(rAngle / 2); + float qz = camera.view.z * sin(rAngle / 2); + float qw = cos(rAngle / 2); + + glm::quat q = glm::quat(qw, qx, qy, qz); + camera.right = q * camera.up; + //set up render camera stuff int arraylen = camera.resolution.x * camera.resolution.y; state.image.resize(arraylen); @@ -176,6 +192,17 @@ int Scene::loadMaterial(string materialid) { newMaterial.emittance = atof(tokens[1].c_str()); } } + + string line; + utilityCore::safeGetline(fp_in, line); + while (!line.empty() && fp_in.good()) { + vector tokens = utilityCore::tokenizeString(line); + if (strcmp(tokens[0].c_str(), "SSS") == 0) { + newMaterial.hasSSS = atof(tokens[1].c_str()); + } + + utilityCore::safeGetline(fp_in, line); + } materials.push_back(newMaterial); return 1; } diff --git a/src/sceneStructs.h b/src/sceneStructs.h index baa2e30..1ff507a 100644 --- a/src/sceneStructs.h +++ b/src/sceneStructs.h @@ -36,6 +36,7 @@ struct Material { float hasRefractive; float indexOfRefraction; float emittance; + float hasSSS; }; struct Camera { @@ -44,6 +45,9 @@ struct Camera { glm::vec3 view; glm::vec3 up; glm::vec2 fov; + glm::vec3 right; + glm::vec3 toGrid; + float dof; }; struct RenderState { @@ -53,3 +57,23 @@ struct RenderState { std::vector image; std::string imageName; }; + +struct PathRay { + glm::vec3 color; + glm::vec3 intersect; + glm::vec3 normal; + Ray ray; + int index; + bool terminate; + int matId; + bool hasIntersect; + bool outside; +}; + +struct is_terminated{ + __host__ __device__ + bool operator()(const PathRay &p) + { + return p.terminate; + } +}; \ No newline at end of file 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..a32d2f3 --- /dev/null +++ b/stream_compaction/efficient.cu @@ -0,0 +1,232 @@ +#include +#include +#include "efficient.h" +#include + +#include +#include +#include +#include +#include +#include +#include + +#define BLOCKSIZE 128 +#define CEILING 7 + +#define FILENAME1 (strrchr(__FILE__, '/') ? strrchr(__FILE__, '/') + 1 : __FILE__) +#define checkCUDAError1(msg) checkCUDAErrorFn1(msg, FILENAME1, __LINE__) +#define ERRORCHECK1 0 +void checkCUDAErrorFn1(const char *msg, const char *file, int line) { +#if ERRORCHECK1 + + cudaDeviceSynchronize(); + 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); +#endif ERRORCHECK1 +} + +namespace StreamCompaction { +namespace Efficient { + + __global__ void filter(int *odata, const PathRay *idata, const int n){ + int k = blockIdx.x*blockDim.x + threadIdx.x; + if (k < n){ + odata[k] = (int)!idata[k].terminate; + } + } + + __global__ void scatter(PathRay *odata, const PathRay *idata, const int *filter, const int *idx, const int n){ + int k = blockIdx.x*blockDim.x + threadIdx.x; + + if (k < n){ + PathRay p = idata[k]; + int f = filter[k]; + int i = idx[k]; + if (f == 1){ + odata[i] = p; + } + } + } + + __global__ void countF(int *c, const int *f, const int *idx, const int n){ + c[0] = f[n - 1] + idx[n - 1]; + } + + /** + * Scan for one single block + */ + __global__ void smallScan(int *small_scan, const int *idata, const int n){ + __shared__ int scanBlock[BLOCKSIZE]; + __shared__ int dArray[CEILING+1]; + int k = blockIdx.x*blockDim.x + threadIdx.x; + + //int ceiling = ilog2ceil(BLOCKSIZE); + + if (threadIdx.x <= CEILING){ + dArray[threadIdx.x] = lround(pow((double)2, (double)(threadIdx.x))); + } + + __syncthreads(); + + if (k < n){ + scanBlock[threadIdx.x] = idata[k]; + } + else { + scanBlock[threadIdx.x] = 0; + } + + // Up sweep + /* + int p1, p2; + for (int d = 0; d < ceiling; d++){ + __syncthreads(); + p1 = lround(pow((double)2, (double)(d + 1))); + p2 = lround(pow((double)2, (double)d)); + if (threadIdx.x % p1 == 0){ + scanBlock[threadIdx.x - 1 + p1] += scanBlock[threadIdx.x - 1 + p2]; + } + } + */ + int p1, p2, tminus = threadIdx.x - 1, d; + for (d = 0; d < CEILING; d++){ + __syncthreads(); + if (threadIdx.x % dArray[d + 1] == 0){ + p2 = tminus + dArray[d]; + p1 = tminus + dArray[d+1]; + scanBlock[p1] += scanBlock[p2]; + } + } + + // Reset root + if (threadIdx.x == 0){ + scanBlock[BLOCKSIZE-1] = 0; + } + + /* + for (int d = ceiling - 1; d >= 0; d--){ + __syncthreads(); + p1 = lround(pow((double)2, (double)(d + 1))); + p2 = lround(pow((double)2, (double)d)); + if (threadIdx.x % p1 == 0){ + int tmp = scanBlock[threadIdx.x - 1 + p2]; + scanBlock[threadIdx.x - 1 + p2] = scanBlock[threadIdx.x - 1 + p1]; + scanBlock[threadIdx.x - 1 + p1] += tmp; + } + } + */ + for (d = CEILING - 1; d >= 0; d--){ + __syncthreads(); + if (threadIdx.x % dArray[d + 1] == 0){ + p2 = tminus + dArray[d]; + p1 = tminus + dArray[d+1]; + int tmp = scanBlock[p2]; + scanBlock[p2] = scanBlock[p1]; + scanBlock[p1] += tmp; + } + } + __syncthreads(); + + small_scan[k] = scanBlock[threadIdx.x]; + } + + __global__ void getBlockTotal(int *block_total, const int *scans, const int * data_in, const int B, const int n){ + int totalIndex = (blockIdx.x + 1)*B - 1; + if (totalIndex < n){ + block_total[blockIdx.x] = scans[totalIndex] + data_in[totalIndex]; + } + else { + block_total[blockIdx.x] = scans[totalIndex]; + } + } + + __global__ void blockIncrement(int *odata, const int *block_incr, const int *small_scan, const int n){ + int k = blockIdx.x*blockDim.x + threadIdx.x; + if (k < n){ + odata[k] = block_incr[blockIdx.x] + small_scan[k]; + } + } + +/** + * Performs prefix-sum (aka scan) on idata, storing the result into odata. + */ +void scan(int n, int *odata, const int *idata) { + int *scan_out; + if (n <= BLOCKSIZE){ + // Base case + cudaMalloc((void**)&scan_out, BLOCKSIZE * sizeof(int)); + // Scan on each block & get total sum + smallScan << <1, BLOCKSIZE >> >(scan_out, idata, n); + cudaMemcpy(odata, scan_out, n * sizeof(int), cudaMemcpyDeviceToHost); + + checkCUDAError1("SC small scan base case"); + } + else { + // Divide into blocks and padding + //int gridSize = lround(ceil((double)n / (double)BLOCKSIZE)); + int gridSize = (n + BLOCKSIZE - 1) / BLOCKSIZE; + + int *block_total; + + // Padding zero + cudaMalloc((void**)&scan_out, gridSize * BLOCKSIZE * sizeof(int)); + cudaMalloc((void**)&block_total, gridSize * sizeof(int)); + + smallScan << > >(scan_out, idata, n); + checkCUDAError1("SC small scan"); + + getBlockTotal << > >(block_total, scan_out, idata, BLOCKSIZE, n); + + // Scan and get block increment; recursively + scan(gridSize, block_total, block_total); + + // Increment block + blockIncrement << > >(odata, block_total, scan_out, n); + checkCUDAError1("SC block increment"); + + cudaFree(block_total); + } + cudaFree(scan_out); +} + +/** + * Performs stream compaction on idata, storing the result into odata. + * All zeroes are discarded. + * + * @param n The number of elements in idata. + * @param idata The array of elements to compact. + * @returns The number of elements remaining after compaction. + */ +void compact(int n, int* f, int* idx, PathRay *dv_out, PathRay *idata, int* c) { + // Padding + //int gridSize = (n + BLOCKSIZE - 1) / BLOCKSIZE; + //int m = (int)pow((double)2, (double)ilog2ceil(gridSize)); + //int tsize = m > gridSize ? m : gridSize; + //int tsize = gridSize; + int tsize = (n + BLOCKSIZE - 1) / BLOCKSIZE; + + // Filter + filter << > >(f, idata, n); + checkCUDAError1("SC filter"); + + // Scan + scan(n, idx, f); + + // Scatter + scatter << > >(dv_out, idata, f, idx, n); + checkCUDAError1("SC scatter"); + + // Get new array size + countF<<<1, 1>>>(c, f, idx, n); +} + +} +} diff --git a/stream_compaction/efficient.h b/stream_compaction/efficient.h new file mode 100644 index 0000000..7b4781a --- /dev/null +++ b/stream_compaction/efficient.h @@ -0,0 +1,22 @@ +#pragma once +#include + +__host__ __device__ inline int ilog2(int x) { + int lg = 0; + while (x >>= 1) { + ++lg; + } + return lg; +} + +__host__ __device__ inline int ilog2ceil(int x) { + return ilog2(x - 1) + 1; +} + +namespace StreamCompaction { +namespace Efficient { + void scan(int n, int *odata, const int *idata); + + void compact(int n, int *f, int *idx, PathRay *dv_out_tmp, PathRay *idata, int *c); +} +}