diff --git a/README.md b/README.md index 93aaa2e..7dfb217 100644 --- a/README.md +++ b/README.md @@ -3,287 +3,19 @@ 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) +* Xinyue Zhu +* Tested on: ( Windows 10, i5-5200U @ 2.20GHz 8GB, GTX 960M +window 10 can be really buggy in this project..crushed my computer again and again. -### (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 -* 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++. - * You can use the triangle intersection function `glm::intersectRayTriangle`. - -This 'extra features' list is not comprehensive. If you have a particular idea -you would like to implement (e.g. acceleration structures, etc.), please -contact us first. - -For each extra feature, you must provide the following analysis: - -* Overview write-up of the feature -* Performance impact of the feature -* If you did something to accelerate the feature, what did you do and why? -* Compare your GPU version of the feature to a HYPOTHETICAL CPU version - (you don't have to implement it!) Does it benefit or suffer from being - implemented on the GPU? -* How might this feature be optimized beyond your current implementation? - -## Base Code Tour - -You'll be working in the following files. Look for important parts of the code: -search for `CHECKITOUT`. You'll have to implement parts labeled with `TODO`. -(But don't let these constrain you - you have free rein!) - -* `src/pathtrace.cu`: path tracing kernels, device functions, and calling code - * `pathtraceInit` initializes the path tracer state - it should copy - scene data (e.g. geometry, materials) from `Scene`. - * `pathtraceFree` frees memory allocated by `pathtraceInit` - * `pathtrace` performs one iteration of the rendering - it handles kernel - launches, memory copies, transferring some data, etc. - * See comments for a low-level path tracing recap. -* `src/intersections.h`: ray intersection functions - * `boxIntersectionTest` and `sphereIntersectionTest`, which take in a ray and - a geometry object and return various properties of the intersection. -* `src/interactions.h`: ray scattering functions - * `calculateRandomDirectionInHemisphere`: a cosine-weighted random direction - in a hemisphere. Needed for implementing diffuse surfaces. - * `scatterRay`: this function should perform all ray scattering, and will - call `calculateRandomDirectionInHemisphere`. See comments for details. -* `src/main.cpp`: you don't need to do anything here, but you can change the - program to save `.hdr` image files, if you want (for postprocessing). - -### Generating random numbers - -``` -thrust::default_random_engine rng(hash(index)); -thrust::uniform_real_distribution u01(0, 1); -float result = u01(rng); -``` - -There is a convenience function for generating a random engine using a -combination of index, iteration, and depth as the seed: - -``` -thrust::default_random_engine rng = random_engine(iter, index, depth); -``` - -### 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 -* SCATTER (float scatter) //scatter flag, 0 for no, 1 for yes -* ABSCOEFF (float r) (float b) (float g) //absorption coefficient for scattering -* RSCTCOEFF (float rsctcoeff) //reduced scattering coefficient -* 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 -* frame (frame number) //indicates start of one frame of parameters - ignore this -* 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 -* frame (frame number) //indicates start of one frame of parameters - ignore this -* 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. +1.refraction & perfect reflection &diffuse color &direct light from ceiling light +|![](img/2.png) +2.without depth of field +|![](img/camera1.png) +3.with depth of field +|![](img/camera2.png) +4.without direct light +|![](img/cornell.10e-4offset.png) diff --git a/img/1.png b/img/1.png new file mode 100644 index 0000000..53e826d Binary files /dev/null and b/img/1.png differ diff --git a/img/2.png b/img/2.png new file mode 100644 index 0000000..9ddd22f Binary files /dev/null and b/img/2.png differ diff --git a/img/camera1.png b/img/camera1.png new file mode 100644 index 0000000..2ca4364 Binary files /dev/null and b/img/camera1.png differ diff --git a/img/camera2.png b/img/camera2.png new file mode 100644 index 0000000..9097173 Binary files /dev/null and b/img/camera2.png differ diff --git a/img/cornell.10e-4offset.png b/img/cornell.10e-4offset.png new file mode 100644 index 0000000..ebc4770 Binary files /dev/null and b/img/cornell.10e-4offset.png differ diff --git a/img/cornell.5000camera_ray+-0.001.png b/img/cornell.5000camera_ray+-0.001.png new file mode 100644 index 0000000..8563f61 Binary files /dev/null and b/img/cornell.5000camera_ray+-0.001.png differ diff --git a/img/cornell5000+10e-5offset.png b/img/cornell5000+10e-5offset.png new file mode 100644 index 0000000..8b5e477 Binary files /dev/null and b/img/cornell5000+10e-5offset.png differ diff --git a/scenes/camera.txt b/scenes/camera.txt new file mode 100644 index 0000000..a50fe6d --- /dev/null +++ b/scenes/camera.txt @@ -0,0 +1,216 @@ +// Emissive material (light) +MATERIAL 0 +RGB 1 1 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 5 + +// Diffuse white +MATERIAL 1 +RGB .95 .95 .95 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 0 +REFRIOR 0 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 0 + +// Diffuse red +MATERIAL 2 +RGB .85 .35 .35 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 0 +REFRIOR 0 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 0 + +// Diffuse green +MATERIAL 3 +RGB .65 .85 .35 +SPECEX 0 +SPECRGB 1 1 1 +REFL 0 +REFR 0 +REFRIOR 0 +SCATTER 0 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 0 + +//refraction,transparent +MATERIAL 4 +RGB .35 .25 .75 +SPECEX 2.0 +SPECRGB 0.2 0.1 0.7 +REFL 1 +REFR 1 +REFRIOR 1.3 +SCATTER 0.7 +ABSCOEFF 0.4 0.4 0.4 +RSCTCOEFF 0 +EMITTANCE 0 + +//perfect specular +MATERIAL 5 +RGB .35 .85 .35 +SPECEX 1.8 +SPECRGB 0.8 0.9 0.8 +REFL 1 +REFR 0 +REFRIOR 1.3 +SCATTER 0.7 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 0 + +// Camera +CAMERA +RES 800 800 +FOVY 45 +ITERATIONS 5000 +DEPTH 8 +FILE cornell +frame 0 +EYE 0 4 7 +VIEW 0 0 -1 +UP 0 1 0 + +TEXTURE + +// Ceiling light +OBJECT 0 +cube +material 0 +frame 0 +TRANS 0 10 0 +ROTAT 0 0 0 +SCALE 3 .3 3 + +// Floor +OBJECT 1 +cube +material 1 +frame 0 +TRANS 0 0 0 +ROTAT 0 0 0 +SCALE 10 .01 10 + +// Ceiling +OBJECT 2 +cube +material 1 +frame 0 +TRANS 0 10 0 +ROTAT 0 0 90 +SCALE .01 10 10 + +// Back wall +OBJECT 3 +cube +material 1 +frame 0 +TRANS 0 5 -5 +ROTAT 0 90 0 +SCALE .01 10 10 + +// Left wall +OBJECT 4 +cube +material 2 +frame 0 +TRANS -5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +// Right wall +OBJECT 5 +cube +material 3 +frame 0 +TRANS 5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +// Sphere/////////////////////////////////////////////// +OBJECT 6 +sphere +material 3 +frame 0 +TRANS 3 3 3 +ROTAT 0 0 0 +SCALE 2 2 2 + +// Sphere +OBJECT 7 +sphere +material 3 +frame 0 +TRANS 3 3 1 +ROTAT 0 0 0 +SCALE 2 2 2 + +// Sphere +OBJECT 8 +sphere +material 3 +frame 0 +TRANS 3 3 -1 +ROTAT 0 0 0 +SCALE 2 2 2 + +// Sphere +OBJECT 9 +sphere +material 3 +frame 0 +TRANS 3 3 -3 +ROTAT 0 0 0 +SCALE 2 2 2 + +OBJECT 10 +sphere +material 3 +frame 0 +TRANS -3 3 3 +ROTAT 0 0 0 +SCALE 2 2 2 + +// Sphere +OBJECT 11 +sphere +material 3 +frame 0 +TRANS -3 3 1 +ROTAT 0 0 0 +SCALE 2 2 2 + +// Sphere +OBJECT 12 +sphere +material 3 +frame 0 +TRANS -3 3 -1 +ROTAT 0 0 0 +SCALE 2 2 2 + +// Sphere +OBJECT 13 +sphere +material 3 +frame 0 +TRANS -3 3 -3 +ROTAT 0 0 0 +SCALE 2 2 2 \ No newline at end of file diff --git a/scenes/cornell.txt b/scenes/cornell.txt index 7c71b63..dbfdfdd 100644 --- a/scenes/cornell.txt +++ b/scenes/cornell.txt @@ -13,7 +13,7 @@ EMITTANCE 5 // Diffuse white MATERIAL 1 -RGB .98 .98 .98 +RGB .95 .95 .95 SPECEX 0 SPECRGB 1 1 1 REFL 0 @@ -50,6 +50,32 @@ ABSCOEFF 0 0 0 RSCTCOEFF 0 EMITTANCE 0 +//refraction,transparent +MATERIAL 4 +RGB .35 .25 .25 +SPECEX 2.0 +SPECRGB 0.2 0.1 0.1 +REFL 1 +REFR 1 +REFRIOR 1.3 +SCATTER 0.7 +ABSCOEFF 0.4 0.4 0.4 +RSCTCOEFF 0 +EMITTANCE 0 + +//perfect specular +MATERIAL 5 +RGB .35 .85 .35 +SPECEX 1.8 +SPECRGB 0.8 0.9 0.8 +REFL 1 +REFR 0 +REFRIOR 1.3 +SCATTER 0.7 +ABSCOEFF 0 0 0 +RSCTCOEFF 0 +EMITTANCE 0 + // Camera CAMERA RES 800 800 @@ -62,6 +88,7 @@ EYE 0.0 5 10.5 VIEW 0 0 -1 UP 0 1 0 +TEXTURE // Ceiling light OBJECT 0 @@ -122,6 +149,25 @@ OBJECT 6 sphere material 1 frame 0 -TRANS -1 4 -1 +TRANS -3 7 -1 +ROTAT 0 0 0 +SCALE 2 2 2 + +// Sphere +OBJECT 7 +sphere +material 4 +frame 0 +TRANS 3 2.5 -1 +ROTAT 0 0 0 +SCALE 4 4 4 + +// Sphere +OBJECT 8 +sphere +material 5 +frame 0 +TRANS -3 2 -1 ROTAT 0 0 0 SCALE 3 3 3 + diff --git a/src/ImageStone.h b/src/ImageStone.h new file mode 100644 index 0000000..d61b547 --- /dev/null +++ b/src/ImageStone.h @@ -0,0 +1,44 @@ +/* + Copyright (C) =USTC= Fu Li + + Author : Fu Li + Create : 1997-9-1 + Home : http://www.phoxo.com + Mail : crazybitwps@hotmail.com + + This file is part of ImageStone + + The code distributed in the hope that it will be useful, + but WITHOUT ANY WARRANTY; without even the implied warranty of + MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. + + Redistribution and use the source code, with or without modification, + must retain the above copyright. +*/ +#ifndef IMAGESTONE_HEADER_9711_H +#define IMAGESTONE_HEADER_9711_H + +#ifdef _MSC_VER + #pragma warning(push) + #pragma warning(disable : 4100) // disable MSVC warning: warning C4100: 'p' : unreferenced formal parameter +#endif + +//------------------------------------------------------------------------------------- +#include "include/image.h" + +#include "include/image_codec/codecfactory_gdiplus.h" + +#ifdef IMAGESTONE_USE_FREEIMAGE + #include "include/image_codec/codecfactory_freeimage.h" +#endif + +#include "include/effect/effect.h" + +#include "include/post_implement.inl" +//------------------------------------------------------------------------------------- + +#ifdef _MSC_VER + #pragma warning(pop) +#endif + +#endif diff --git a/src/interactions.h b/src/interactions.h index 22c1710..13a7cbe 100644 --- a/src/interactions.h +++ b/src/interactions.h @@ -9,36 +9,38 @@ */ __host__ __device__ glm::vec3 calculateRandomDirectionInHemisphere( - glm::vec3 normal, thrust::default_random_engine &rng) { - thrust::uniform_real_distribution u01(0, 1); +glm::vec3 normal, thrust::default_random_engine &rng) {//generate a random number + thrust::uniform_real_distribution u01(0, 1); - float up = sqrt(u01(rng)); // cos(theta) - float over = sqrt(1 - up * up); // sin(theta) - float around = u01(rng) * TWO_PI; + float up = sqrt(u01(rng)); // cos(theta) + float over = sqrt(1 - up * up); // sin(theta) + float around = u01(rng) * TWO_PI; - // 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. + // 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); - } + 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)); + // 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; + return up * normal + + cos(around) * over * perpendicularDirection1 + + sin(around) * over * perpendicularDirection2; } /** @@ -47,22 +49,86 @@ glm::vec3 calculateRandomDirectionInHemisphere( * 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: + * - 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. + *//*float calculateFresnelReflectionCoefficient(glm::vec3 direction, glm::vec3 normal, float indexOfRefraction) { + 49 float r0 = glm::pow((1.0f - indexOfRefraction) / (1.0f + indexOfRefraction), 2); + 50 return r0 + (1.0f - r0) * glm::pow(1.0f - glm::dot(normal, -direction), 5); + 51 + }*/ + +__host__ __device__ float R_coeff(float theta, float n2){ + //air :n1=1; ice:1.3 + + float n1 = 1.0; //n2 = 1.3; + float R0 = pow((n1 - n2) / (n1 + n2), 2); + float Rtheta = R0 + (1 - R0)*pow((1 - theta), 5); + return Rtheta; } +__host__ __device__ void reflectRay(Ray &r, glm::vec3 normal, glm::vec3 intersection, Material m){ + r.direction = glm::reflect(r.direction, normal); + r.origin = intersection; + r.hitcolor *= m.specular.color; +} + +__host__ __device__ +void scatterRay(Ray &ray, glm::vec3 &color, glm::vec3 intersect, glm::vec3 normal, bool outside, glm::vec3 emittedColor, +const Material &m, thrust::default_random_engine &rng) { + + if (m.hasRefractive){ + + float theta = glm::dot(-ray.direction, normal); + float Rtheta = R_coeff(theta, m.indexOfRefraction); + thrust::uniform_real_distribution u01(0, 1); + glm::vec3 refract_ray, reflect_ray; + if (u01(rng) <1-Rtheta)//refract + { + //refractRay(ray, normal, intersect, m, outside); + float n = 0; + if (outside) { + //float + float n = 1.0f / m.indexOfRefraction; // why I have to wirte the float agian??? + } + + float sintheta2 = glm::pow(n, 2) * glm::pow(glm::dot(-normal, ray.direction), 2); + if (1.0-sintheta2>=0){ + glm::vec3 r1 = ray.direction; + glm::vec3 r2 = n*r1 - (n*glm::dot(r1, normal) + float(glm::sqrt(1.0-sintheta2)))*normal; + ray.direction = r2; + ray.origin = intersect + ray.direction * 0.0005f;//maybe because of the shadow acne it canot shows itselt... + //so I add an offset but why there is the black edge... + } + //else I assume they are absored + + } + else{ + reflectRay(ray, normal, intersect, m); + } + } + //} + if (m.hasReflective&&(!m.hasRefractive)){ + if (m.specular.exponent > 1){ + reflectRay(ray, normal, intersect, m);//perfect mirror + } + } + + //diffuse + if ((!m.hasReflective) && (!m.hasRefractive)){ + glm::vec3 diffuse_ray = calculateRandomDirectionInHemisphere(normal, rng); + ray.direction = glm::normalize(diffuse_ray); + ray.origin = intersect; + ray.hitcolor *= m.color; + } + //else color + } diff --git a/src/intersections.h b/src/intersections.h index 5f3613d..dcaae11 100644 --- a/src/intersections.h +++ b/src/intersections.h @@ -1,144 +1,165 @@ #pragma once #include - +#include #include "sceneStructs.h" #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. - */ -__host__ __device__ glm::vec3 getPointOnRay(Ray r, float t) { - return r.origin + (t - .0001f) * glm::normalize(r.direction); +* 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); } /** - * 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); +* 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); } // 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__ void motionBlur(Geom geo, glm::vec3 translate, int frame){ + geo.translation += translate*float(frame)*0.01f; + geo.transform = utilityCore::buildTransformationMatrix(geo.translation, geo.rotation, geo.scale); + geo.inverseTransform = glm::inverse(geo.transform); + geo.invTranspose = glm::inverseTranspose(geo.transform);//hpp +}*/ + __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, 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. - */ +* 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, bool &outside,bool go=0,int frame =1,glm::vec3& mostiontranslate=glm::vec3(0,0,0)) { +// if (go) + // motionBlur(sphere, mostiontranslate, frame); + 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.0; + } + 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); + } diff --git a/src/main.cpp b/src/main.cpp index 77671f4..0b264ba 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -29,7 +29,7 @@ int main(int argc, char** argv) { const char *sceneFile = argv[1]; // Load scene file - scene = new Scene(sceneFile); + scene = new Scene(sceneFile);//marerial, object, camera // Set up camera stuff from loaded path tracer settings iteration = 0; @@ -41,8 +41,8 @@ int main(int argc, char** argv) { init(); // GLFW main loop - mainLoop(); - + mainLoop();//runCUDA +// getchar(); return 0; } @@ -87,7 +87,7 @@ void runCuda() { // Map OpenGL buffer object for writing from CUDA on a single GPU // No data is moved (Win & Linux). When mapped to CUDA, OpenGL should not use this buffer - + if (iteration == 0) { pathtraceFree(); pathtraceInit(scene); diff --git a/src/pathtrace.cu b/src/pathtrace.cu index 68721e6..940390c 100644 --- a/src/pathtrace.cu +++ b/src/pathtrace.cu @@ -13,152 +13,630 @@ #include "pathtrace.h" #include "intersections.h" #include "interactions.h" +#include +#define ERRORCHECK 1 #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; - } +#if ERRORCHECK + 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); + fprintf(stderr, "CUDA error"); + if (file) { + fprintf(stderr, " (%s:%d)", file, line); + } + fprintf(stderr, ": %s: %s\n", msg, cudaGetErrorString(err)); +# ifdef _WIN32 + getchar(); +# endif + exit(EXIT_FAILURE); +# endif } -__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)); + +__host__ __device__ +thrust::default_random_engine makeSeededRandomEngine(int iter, int index, int depth) { + int h = utilhash((1 << 31) | (depth << 22) | 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; + 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]; + 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); + 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; - } + // 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 = NULL; static glm::vec3 *dev_image = NULL; // TODO: static variables for device memory, scene/camera info, etc // ... +static Ray *dev_ray = NULL; +static int *dev_bool = NULL; +static Ray *dev_compacted = NULL; +static Material *dev_m; +static Geom *dev_geo = NULL; +static Geom *dev_geoms; +static glm::vec3 *dev_test = NULL; +static glm::vec3 *dev_colormap = NULL; +static int * dev_combool = NULL; +static int * dev_resultint = NULL; +static Ray * dev_resultray = NULL; +static Ray *camera_ray = NULL; +static Ray *dev_compatedRay = NULL; +static bool* dev_terminate = NULL; +static int * dev_newcombinedSumdata = NULL; +static Ray *dev_compactResult=NULL; + +__device__ void depth_of_field(Camera cmr,int iter, int index,Ray &ray,int x,int y){ + glm::vec3 horizontal, middle, vertical; + + glm::vec3 F = glm::normalize(cmr.view); + glm::vec3 R = glm::normalize(glm::cross(F, cmr.up)); + glm::vec3 U = glm::normalize(glm::cross(R, F)); + float len = glm::length(cmr.view); + int width = cmr.resolution.x; + int height = cmr.resolution.y; + float alpha = cmr.fov.y*PI / 180.f; + glm::vec3 V = U*len *tan(alpha); + float temp = width*1.0 / (height*1.0); + glm::vec3 H = temp*glm::length(F)*R; + //camera componet(EYE:0,4,7) + float lens_radius = 2.0f; + float focal_distance = 2.f;//aperon focus + // step1.Sample a random point on the lense + thrust::default_random_engine rng = makeSeededRandomEngine(iter, index, 1); + thrust::uniform_real_distribution u01(-1, 1); + //a concentric sample disk + float theta = PI * u01(rng); + float dx = lens_radius*cosf(theta); + float dy = lens_radius * sinf(theta); + glm::vec2 point = glm::vec2(dx, dy); + glm::vec3 position = cmr.position + glm::vec3(dx, dy, 0); + float ft = focal_distance / cmr.view.z;//cmr.position.z; + thrust::uniform_real_distribution uee(-EPSILON, EPSILON); + + float offset = uee(rng); + float xx, yy; + xx = 2.0* x / width - 1.0 + offset; + yy = 1.0 - 2.0* y / height + offset; + + glm::vec3 point_pos = cmr.view + position + xx*H + yy*V; + //glm::vec3 point_focus = cmr.view*ft + position; + //ray.origin = cmr.position + (dx * R + dy * U);; + ray.origin = position; + ray.direction = glm::normalize(point_pos - ray.origin); + + + +} +__global__ void SetDevRay(Camera cmr, Ray *dev_ray, int iter){ + + int x = (blockIdx.x * blockDim.x) + threadIdx.x; + int y = (blockIdx.y * blockDim.y) + threadIdx.y; + + if (x < cmr.resolution.x && y < cmr.resolution.y) { + int index = x + (y * cmr.resolution.x); + glm::vec3 F = glm::normalize(cmr.view); + glm::vec3 R = glm::normalize(glm::cross(F, cmr.up)); + glm::vec3 U = glm::normalize(glm::cross(R, F)); + float len = glm::length(cmr.view); + int width = cmr.resolution.x; + int height = cmr.resolution.y; + float alpha = cmr.fov.y*PI / 180.f; + glm::vec3 V = U*len *tan(alpha); + float temp = width*1.0 / (height*1.0); + glm::vec3 H = temp*glm::length(F)*R; + float xx; + float yy; + //jittering rays within an aperture + + + thrust::default_random_engine rng = makeSeededRandomEngine(iter, index, 1); + thrust::uniform_real_distribution uee(-EPSILON, EPSILON); + + float offset = uee(rng); + + xx = 2.0* x / width - 1.0 + offset; + yy = 1.0 - 2.0* y / height + offset; + + glm::vec3 point_pos = cmr.view + cmr.position + xx*H + yy*V;//glm::vec3 M_ = cmr.position + C_; + //screen point to world point + dev_ray[index].direction = glm::normalize(point_pos - cmr.position ); + dev_ray[index].origin = cmr.position; + dev_ray[index].terminate = false; + dev_ray[index].hitcolor = WHITE; + ///////Camera cmr,int iter, int index,Ray &ray,int x,int y + depth_of_field(cmr, iter, index, dev_ray[index],x,y); + } +} 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; + 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_compactResult, pixelcount * sizeof(Ray)); + cudaMemset(dev_compactResult, 0, pixelcount * sizeof(Ray)); + + int geoSize = hst_scene->geoms.size()*sizeof(Geom); + Geom * hst_geoms = (Geom *)malloc(geoSize); + std::copy(hst_scene->geoms.begin(), hst_scene->geoms.end(), hst_geoms); + cudaMalloc((void**)&dev_geoms, geoSize); + cudaMemcpy(dev_geoms, hst_geoms, geoSize, cudaMemcpyHostToDevice); + + ////////////// + /////// + //camera_ray = new Ray[pixelcount]; + + int nG = hst_scene->geoms.size(); + int nM = hst_scene->materials.size(); + + Material * mm = new Material[nM]; + Geom *gg = new Geom[nG]; + for (int i = 0; i < nM; i++){ + mm[i] = hst_scene->materials[i]; + } + for (int i = 0; i < nG; i++){ + gg[i] = hst_scene->geoms[i]; + } + cudaMalloc((void**)&dev_m, nM*sizeof(Material));//n*sizeof(M) + cudaMemcpy(dev_m, mm, nM*sizeof(Material), cudaMemcpyHostToDevice); + + cudaMalloc((void**)&dev_geo, nG*sizeof(Geom)); + cudaMemcpy(dev_geo, gg, nG*sizeof(Geom), cudaMemcpyHostToDevice); + + // TODO: initialize the above static variables added above + cudaMalloc(&dev_ray, pixelcount * sizeof(Ray)); + cudaMemset(dev_ray, 0, pixelcount * sizeof(Ray)); - 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_colormap, pixelcount * sizeof(glm::vec3)); + cudaMemset(dev_colormap, 0, pixelcount * sizeof(glm::vec3)); + + checkCUDAError("pathtraceInit"); - checkCUDAError("pathtraceInit"); } -void pathtraceFree() { - cudaFree(dev_image); // no-op if dev_image is null - // TODO: clean up the above static variables +__global__ void generateIamge(Camera cam, int iter, glm::vec3 *image, Ray *r) { + + /*int id = (blockIdx.x * blockDim.x) + threadIdx.x; + //if (id < (cam.resolution.x * cam.resolution.y)) { + //if ()*/ + /*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 id = x + (y * cam.resolution.x); + if (r[id].terminate) + image[id] += r[id].hitcolor; + }*/ + int id = (blockIdx.x * blockDim.x) + threadIdx.x; + + if (id < cam.resolution.x *cam.resolution.y) { + + if (r[id].terminate){ + image[id] += r[id].hitcolor; + //r[id].hitcolor = glm::vec3(); + } + } +} +__device__ void directlightcheking(Ray &r, Geom *dev_geom, int nG, const int light, Material m_light, int iter, int id){ + Ray light_ray; + //find random place + glm::vec3 point; + float result[3]; + if (dev_geom[light].type == CUBE){ + thrust::default_random_engine rng = makeSeededRandomEngine(iter, id, 1); + thrust::uniform_real_distribution side(0, 6); + thrust::uniform_real_distribution x01(-0.5, 0.5); + thrust::uniform_real_distribution y01(-0.5, 0.5); + int s = (int)side(rng); + int c = s % 3; + result[c] = s > 2.0 ? 1.f : 0.f; + result[(c + 1) % 3] = x01(rng); + result[(c + 2) % 3] = y01(rng); + point = glm::vec3(result[0], result[1], result[2]); + + glm::vec4 pt = glm::vec4(point, 1)* dev_geom[light].transform; + point = glm::vec3(pt[0], pt[1], pt[2]); + } + + if (dev_geom[light].type == SPHERE){ + thrust::default_random_engine rng = makeSeededRandomEngine(iter, id, 1); + thrust::uniform_real_distribution u01(0, 1); + thrust::uniform_real_distribution v01(0, 1); + + //float radius = dev_geom[light].scale + float u = u01(rng); + float v = v01(rng); + float theta = 2 * PI * u; + float phi = acos(2 * v - 1); + float x = 0.5 * sin(phi) * cos(theta); + float y = 0.5 * sin(phi) * sin(theta); + float z = 0.5 * cos(phi); - checkCUDAError("pathtraceFree"); + point = glm::vec3(x, y, z); + point *= dev_geom[light].translation; + } + + light_ray.direction = point - r.origin; + light_ray.origin = r.origin; + float t; + glm::vec3 intersectionPoint; + glm::vec3 normal; + bool outside; + for (int i = 0; i < nG; i++){ + if (i != light){ + if (dev_geom[i].type == SPHERE){ + t = sphereIntersectionTest(dev_geom[i], r, intersectionPoint, normal, outside); + } + if (dev_geom[i].type == CUBE){ + t = boxIntersectionTest(dev_geom[i], r, intersectionPoint, normal, outside); + } + } + if (i == light)t = -1; + } + + if (t < 0){ + r.hitcolor *= m_light.emittance*m_light.color; + } } +__global__ void raytracing(int frame, Ray *r, int CurrentRayNumber, Geom *dev_geom, int nG, int nM, Material *dev_m, int iter, int traced, int currentd) +{ + //glm::vec3 s_translate = glm::vec3(1, 0, 0); + int id = (blockIdx.x * blockDim.x) + threadIdx.x; + if (id < CurrentRayNumber){ + if (!r[id].terminate){ + int go_sphere = 2; + glm::vec3 normal; + glm::vec3 intersectionPoint; + bool outside = false; + float t = -2.0; + int mId = 0; + float mint = 30000; + bool mark = false; + Material m_light; + int light; + glm::vec3 F_normal; + glm::vec3 F_intersectionPoint; + bool F_outside = false; + float F_t = -1.0; + for (int i = 0; i < nM; i++){ + if (dev_m[i].emittance>0)//light + { + light = i; + m_light = dev_m[i]; + } + } + + for (int i = 0; i < nG; i++){ + int ll = 0; + if (dev_geom[i].type == SPHERE){ + ll++; + if (ll == go_sphere){ + t = sphereIntersectionTest(dev_geom[i], r[id], intersectionPoint, normal, outside, 0, frame); + } + else + t = sphereIntersectionTest(dev_geom[i], r[id], intersectionPoint, normal, outside, 0, frame); + } + if (dev_geom[i].type == CUBE){ + t = boxIntersectionTest(dev_geom[i], r[id], intersectionPoint, normal, outside); + } + if (t<0)continue; + if (t > 0 && t < mint){//if I want to find the nearest intersect object + mId = dev_geom[i].materialid; + mint = t; + mark = true; + F_intersectionPoint = intersectionPoint; + F_normal = normal; + F_outside = outside; + F_t = t; -/** - * 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; + thrust::default_random_engine rng = makeSeededRandomEngine(iter, id, currentd); + thrust::uniform_real_distribution uee(0, EPSILON * 10); + float offset = uee(rng); + //along the normal + F_intersectionPoint += offset*F_normal; + } + } + if (F_t < 0)//hit nothing terminate + { + r[id].terminate = true; + r[id].hitcolor *= BLACK; + + } + else if (F_t > 0)//hit something + { + glm::vec3 emmited_c; + glm::vec3 color; + Material m = dev_m[mId];//if (mId<0 || mId>nM):erro + if (m.emittance > 0){//***if hit light + r[id].terminate = true; + r[id].hitcolor *= m.emittance*m.color; + } + else{ + if (currentd == traced - 1){ + r[id].terminate = true; + directlightcheking(r[id], dev_geom, nG, light, m_light, iter, id); + } + thrust::default_random_engine rng = makeSeededRandomEngine(iter, id, currentd); + scatterRay(r[id], color, F_intersectionPoint, F_normal, outside, emmited_c, m, rng); + //r[id].hitcolor *= color; + } + } + } + } +} - if (x < cam.resolution.x && y < cam.resolution.y) { - int index = x + (y * cam.resolution.x); +__global__ void MapBool(int *dev_bool, Ray *dev_ray, int n) { + + int id = (blockIdx.x * blockDim.x) + threadIdx.x; + if (id < n){ + if (dev_ray[id].terminate)dev_bool[id] = 0; + else dev_bool[id] = 1; + } +} - 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 kernScatter(int n, Ray *odata, Ray *idata, const int *bools, const int *indices) +{ + int k = (blockIdx.x * blockDim.x) + threadIdx.x; + if (k < n){ + if (bools[k] ){ + int t = indices[k]; + odata[t] = idata[k]; + } + } } +__global__ void EfficientScan(int n, int *outdata, const int *indata){ +//current size + //Example 39-2 in Gems + extern __shared__ int temp[];//stores the updated bool + + //int id = (blockIdx.x * blockDim.x) + threadIdx.x;//tid + int thid = threadIdx.x; + int offset = 1; + int blocks = blockIdx.x * blockDim.x * 2; + + temp[2 * thid] = indata[2 * thid + blocks]; + temp[2 * thid + 1] = indata[2 * thid + blocks * 2 + 1]; + //...do in block maybe.. + // int n = blockDim.x * 2; + for (int d = n >> 1; d > 0; d >>= 1) //d = 0 to log2 n – 1 + { + __syncthreads(); + if (thid < d){ + int ai = offset*(2 * thid + 1) - 1; + int bi = offset*(2 * thid + 2) - 1; + temp[bi] += temp[ai]; + } + offset *= 2; + } + if (thid == 0) //clear the last element + temp[n - 1] = 0; + + for (int d = 1; d < n; d *= 2) + { + offset >>= 1; + __syncthreads(); + + if (thid < d) { + int ai = offset*(2 * thid + 1) - 1; + int bi = offset*(2 * thid + 2) - 1; + int t = temp[ai]; + temp[ai] = temp[bi]; + temp[bi] += t; + } + } + __syncthreads(); + + outdata[2 * thid + blocks] = indata[2 * thid]; + outdata[2 * thid + blocks + 1] = indata[2 * thid + 1]; +} + +__global__ void BlockSums(int n, int *odata, const int *idata) { + int index = (blockIdx.x * blockDim.x) + threadIdx.x; + odata[index] = idata[(index + 1) * 128 - 1]; +} + + + +__global__ void BlcockIncrement(int n, int *dev_data, const int *increments) { + int id = (blockIdx.x * blockDim.x) + threadIdx.x; + //add back to each block + if (id < n){ + // new[0],new[1]...new[127] =new[0],new[1]...new[127]+increment[0]; + // new[128],new[129]...new[255] =new[128],new[129]..new[255]+increment[1]; + dev_data[id] += increments[blockIdx.x]; + } +} + + + + +int StreamCompact(Ray *dev_ray, Ray *result ,int raynumber){ + + //step1.compute temp bool ray + cudaMalloc((void**)&dev_bool, raynumber * sizeof(int)); + cudaMemset(dev_bool, 0, raynumber * sizeof(int)); + cudaMalloc((void**)&dev_combool, raynumber * sizeof(int)); + cudaMemset(dev_combool, 0, raynumber * sizeof(int)); + dim3 GridSize = (raynumber + 128 - 1) / 128; + dim3 BlockSize = 128; + int *s=new int[raynumber]; + s[0] = -1; + + //const dim3 blockSize2d(8, 8); +// const dim3 blocksPerGrid2d(64) + +// <> +// mapBool(int *dev_bool, Ray *dev_ray, int n) + MapBool <<< GridSize, BlockSize >> >(dev_bool, dev_ray, raynumber); + cudaMemcpy(s, dev_bool, raynumber*sizeof(int), cudaMemcpyDeviceToHost); + //checkCUDAError("mapbool"); + + //************************************/ + //step2.scan Scan(raynumber, dev_compactbool, dev_bool); + dim3 dim3_p0 = GridSize; + int int_p0 = raynumber + 128 - 1 / 128; + cudaMalloc(&dev_newcombinedSumdata, sizeof(int)* 128); + while (int_p0 >= 1){ + int n = (int_p0 + 128 - 1) / 128; + dim3 scanGridSize = (int_p0 + 128 - 1) / 128; + dim3 scanBlockSize = 128;// (?) + + //cudaMalloc(&dev_result, sizeof(int)* 128); + //EfficientScan(int * indata, int *outdata, int n) + EfficientScan << > >(int_p0,dev_combool, dev_bool); + checkCUDAError("scan"); + //step4.write total sum of each block to a new array. + //return the sum of each block + BlockSums << > >(n, dev_newcombinedSumdata,dev_combool); + checkCUDAError("sum"); + //dim3 gBlockSum = p0; + int int_p0 = int_p0 / 128; + + checkCUDAError("blockIncrement"); + } + + BlcockIncrement << < GridSize, 128 >> >(raynumber, dev_combool, dev_combool); + //step3.scatter:kernScatter(int n, Ray *odata, const Ray *idata, const int *bools, const int *indices) + kernScatter << > >(raynumber, result, dev_ray, dev_bool, dev_combool); + checkCUDAError("kernScatter"); + + int current_ray_number; + cudaMemcpy(¤t_ray_number,&(dev_combool[raynumber-1]), sizeof(int), cudaMemcpyDeviceToHost); + cudaMemcpy(dev_ray, dev_resultray, sizeof(Ray)*current_ray_number, cudaMemcpyDeviceToDevice); + + cudaFree(dev_bool); + cudaFree(dev_combool); + checkCUDAError("whatever"); + return current_ray_number; + + +} + + + -/** - * 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. - // * Each path ray is a (ray, color) pair, where color starts as the - // multiplicative identity, white = (1, 1, 1). - // * For debugging, you can output your ray directions as colors. - // * For each depth: - // * Compute one new (ray, color) pair along each path - note - // that many rays will terminate by hitting a light or nothing at all. - // You'll have to decide how to represent your path rays and how - // you'll mark terminated rays. - // * Color is attenuated (multiplied) by reflections off of any object - // surface. - // * You can debug your ray-scene intersections by displaying various - // values as colors, e.g., the first surface normal, the first bounced - // ray direction, the first unlit material color, etc. - // * Add all of the terminated rays' results into the appropriate pixels. - // * Stream compact away all of the terminated paths. - // You may use either 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 - - generateNoiseDeleteMe<<>>(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"); + const int traceDepth = hst_scene->state.traceDepth; + const Camera &cam = hst_scene->state.camera; + const int pixelcount = cam.resolution.x * cam.resolution.y; + + + const dim3 blockSize2d(8, 8); + const dim3 blocksPerGrid2d( + (cam.resolution.x + blockSize2d.x - 1) / blockSize2d.x, + (cam.resolution.y + blockSize2d.y - 1) / blockSize2d.y); + + const int &nG = hst_scene->geoms.size(); + const int &nM = hst_scene->materials.size(); + + /////////////////////////////////////////////////////////////////////////// + + // Recap: + // * Initialize array of path rays (using rays that come out of the camera) + // * You can pass the Camera object to that kernel. + // * Each path ray is a (ray, color) pair, where color starts as the + // multiplicative identity, white = (1, 1, 1). + // * For debugging, you can output your ray directions as colors. + // * For each depth: + // * Compute one new (ray, color) pair along each path (using scatterRay). + // Note that many rays will terminate by hitting a light or hitting + // nothing at all. You'll have to decide how to represent your path rays + // and how you'll mark terminated rays. + // * Color is attenuated (multiplied) by reflections off of any object + // surface. + // * You can debug your ray-scene intersections by displaying various + // values as colors, e.g., the first surface normal, the first bounced + // ray direction, the first unlit material color, etc. + // * Add all of the terminated rays' results into the appropriate pixels. + // * Stream compact away all of the terminated paths. + // You may use either your implementation or `thrust::remove_if` or its + // cousins. + // * Note that you can't really use a 2D kernel launch any more - switch + // to 1D. + // * 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 + //}//use 256 warp erro + + + //SetDevRay << > >(cam, dev_ray, iter); + const dim3 Grid_256 = (pixelcount + 256 - 1) / 256; + const dim3 B_256 = 256; + const dim3 Grid_128 = (pixelcount + 128 - 1) / 128; + const dim3 B_128 = 128; + int current_ray = pixelcount; + int test; + cudaMalloc(&dev_ray, current_ray*sizeof(Ray)); + for (int i = 0; i < traceDepth; i++) + {//remove the ray number from pixel count,so the Grid size needs to be changed + dim3 CurGrid_128 = (current_ray + 128 - 1) / 128; + if (!i) + { + current_ray = pixelcount; + //cudaMalloc(&dev_ray, current_ray*sizeof(Ray)); + SetDevRay << > >(cam, dev_ray, iter);//shoot ray first time&one time + } + //(Ray *r, int CurrentRayNumber, Geom *dev_geom, int nG, int nM, Material *dev_m, int iter, int traced, int currentd) + raytracing << > >(frame, dev_ray, current_ray, dev_geo, nG, nM, dev_m, iter, traceDepth, i); + //checkCUDAError("raytrace"); + //current_ray = StreamCompact(dev_ray, dev_compactResult, current_ray); + //test=StreamCompact(dev_ray, dev_compactResult, current_ray); + //cudaMemcpy(dev_ray, dev_compactResult, sizeof(Ray)*current_ray , cudaMemcpyDeviceToDevice); + } +// printf("%d", test); + generateIamge << > >(cam, iter, dev_image, dev_ray); + cudaFree(dev_ray); + /////////////////////////////////////////////////////////////////////////// + + // 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 pathtraceFree() { + cudaFree(dev_image); // no-op if dev_image is null + // TODO: clean up the above static variables + cudaFree(dev_ray); + cudaFree(dev_compacted); + cudaFree(dev_compactResult); + cudaFree(dev_m); + cudaFree(dev_geo); + cudaFree(dev_colormap); + + checkCUDAError("pathtraceFree"); } diff --git a/src/preview.cpp b/src/preview.cpp index 943456f..bf32749 100644 --- a/src/preview.cpp +++ b/src/preview.cpp @@ -185,6 +185,8 @@ void mainLoop() { glDrawElements(GL_TRIANGLES, 6, GL_UNSIGNED_SHORT, 0); glfwSwapBuffers(window); } + glfwDestroyWindow(window); glfwTerminate(); + } diff --git a/src/scene.cpp b/src/scene.cpp index 0b48f70..429b3c9 100644 --- a/src/scene.cpp +++ b/src/scene.cpp @@ -28,6 +28,7 @@ Scene::Scene(string filename) { loadCamera(); cout << " " << endl; } + } } } @@ -118,7 +119,11 @@ int Scene::loadCamera() { state.traceDepth = atoi(tokens[1].c_str()); } else if (strcmp(tokens[0].c_str(), "FILE") == 0) { state.imageName = tokens[1]; + } + //loadTexture(); + + } string line; @@ -157,6 +162,28 @@ int Scene::loadCamera() { cout << "Loaded camera!" << endl; return 1; } +int Scene::loadTexture(){ + + std::ifstream image; + image.open("D:\\156.jpg", std::ios_base::binary); + + image.seekg(0, ios::end); + int n = image.tellg(); + image.seekg(0, ios::beg); + cout << "n:" << n; + char* res = new char[n]; + glm::vec2 s(1,1); + for (int i = 0; i < n; i++) + res[i] = '5'; + + + bool bit = image.eof(); + + image.read(res, n); + for (int i = 0; i < 100; i++) + printf("%c ", res[i]); + return 1; + } int Scene::loadMaterial(string materialid) { int id = atoi(materialid.c_str()); diff --git a/src/scene.h b/src/scene.h index f29a917..420e886 100644 --- a/src/scene.h +++ b/src/scene.h @@ -16,11 +16,13 @@ class Scene { int loadMaterial(string materialid); int loadGeom(string objectid); int loadCamera(); + int loadTexture(); public: Scene(string filename); ~Scene(); std::vector geoms; std::vector materials; + std::vectortexture; RenderState state; }; diff --git a/src/sceneStructs.h b/src/sceneStructs.h index baa2e30..f9f8a83 100644 --- a/src/sceneStructs.h +++ b/src/sceneStructs.h @@ -13,6 +13,8 @@ enum GeomType { struct Ray { glm::vec3 origin; glm::vec3 direction; + glm::vec3 hitcolor; + bool terminate; }; struct Geom { @@ -25,7 +27,11 @@ struct Geom { glm::mat4 inverseTransform; glm::mat4 invTranspose; }; - +struct Textures{ + glm::vec3 color; + int width; + int height; +}; struct Material { glm::vec3 color; struct { diff --git a/src/utilities.cpp b/src/utilities.cpp index 9c06c68..37cc06f 100644 --- a/src/utilities.cpp +++ b/src/utilities.cpp @@ -20,7 +20,17 @@ float utilityCore::clamp(float f, float min, float max) { return f; } } +int utilityCore::ilog2(int x) { + int lg = 0; + while (x >>= 1) { + ++lg; + } + return lg; +} +int utilityCore::ilog2ceil(int x) { + return ilog2(x - 1) + 1; +} bool utilityCore::replaceString(std::string& str, const std::string& from, const std::string& to) { size_t start_pos = str.find(from); if (start_pos == std::string::npos) diff --git a/src/utilities.h b/src/utilities.h index abb4f27..688040a 100644 --- a/src/utilities.h +++ b/src/utilities.h @@ -13,8 +13,13 @@ #define TWO_PI 6.2831853071795864769252867665590057683943f #define SQRT_OF_ONE_THIRD 0.5773502691896257645091487805019574556476f #define EPSILON 0.00001f - +#define WHITE glm::vec3(1,1,1) +#define BLACK glm::vec3(0,0,0) +#define RED glm::vec3(1,0,0) +#define GREEN glm::vec3(0,1,0) namespace utilityCore { + extern int ilog2(int x); + extern int ilog2ceil(int x); extern float clamp(float f, float min, float max); extern bool replaceString(std::string& str, const std::string& from, const std::string& to); extern glm::vec3 clampRGB(glm::vec3 color);