diff --git a/README.md b/README.md index 110697c..907ccf0 100644 --- a/README.md +++ b/README.md @@ -3,11 +3,79 @@ 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) +* Liang Peng +* Tested on: Windows 10, i7-6700HQ @ 2.6GHz 2.6GHz 8GB, GTX 960M (Personal Laptop) -### (TODO: Your README) +### Feature +* [x] Ray Scattering + * [x] Diffuse + * [x] Refraction + * [x] Specular Reflection + * [x] Glossy Reflection +* [x] Depth of Field +* [x] Stratified Antialiasing +* [x] Performance Analysis -*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. +### Overview +
+The image is rendered by a path tracer on GPU, with trace depth 8 and 5000 samples per pixel. Features include diffuse, glossy, specular reflection, refraction, depth of field, and stratified antialiasing. The image is produced by averaging the results of all large amount of iterations. In each iteration, the camera will shoot a ray through each pixel and intersect with the scene objects. Based on the material at the intersection point, each ray will change color and be reflected towards certain direction and continue tracing. As a ray reaches it trace depth, its color will be accumulated to compute the final color. +### Ray Scattering +* Lambert Diffuse + +Trace Depth 2, SPP 1000 | Trace Depth 8, SPP 1000 +--- | --- + | +_observation_ Images rendered with deeper trace level appear darker, which is due to the calculation of ray color. The resulted color is calculated by multiplying the material color of all intersections along the path. + +* Refraction + +Refraction 1, SPP 1000 | Refraction 2, SPP 5000 +--- | --- + | +_observation_ Since large amount of paths are traced, caustics produced by refraction are free. + +* Specular Reflection + +Trace Depth 2, SPP 1000 | Trace Depth 8, SPP 1000 +--- | --- + | + +* Glossy Reflection + +Less Glossy Floor, SPP 3000 | Glossier Floor, SPP 3000 +--- | --- + | +_obervation_ For different types of material, glossiness can be adjusted by tuning the specular exponent. Larger exponent produces reflection closer to mirror. + +### Depth of Field + +Depth of Field, Far Focal | Depth of Field, Near Focal +--- | --- + | +_obervation_ There are 2 parameters we can play around with in terms of depth of field, lens radius and focal distance. Larger lens radius produces stronger DoF effect. Focal distance affects the distance the virtual camera is focusing on. + +### Antialiasing + +Antialiasing OFF | Antialiasing ON +--- | --- + | +_obersvation_ By jittering the camera ray within the ray's own pixel and performing large number of iterations, the final color of a pixel is computed by averaging all rays going through the pixel, producing antialiased image. + +### Performance Analysis +* Trace Depth + +Depth | 1 | 2 | 3 +:---:|:---:|:---:|:---: +Time | 38 ms | 70 ms | 93 ms +Image | ![](img/depth1.png) | ![](img/depth2.png) | ![](img/depth3.png) +__Depth__ | __4__ | __5__ | __6__ +Time | 110 ms | 120 ms | 140 ms +Image | ![](img/depth4.png) | ![](img/depth5.png) | ![](img/depth6.png) +__Depth__ | __7__ | __8__ | __9__ +Tim3 | 148 ms | 156 ms | 160 ms +Image | ![](img/depth7.png) | ![](img/depth8.png) | ![](img/depth9.png) + +![](img/time_vs_depth.png) + +_Observation_ Apparently, as trace depth increases, image quality becomes better. Interestingly, time spent on 1 iteration increases with decreasing rate, that is probably due to the fact that as a ray traces deeper, it is more likely to hit nothing or terminated, thus need no more calculation. diff --git a/img/DoF2_depth8_3000spp.png b/img/DoF2_depth8_3000spp.png new file mode 100644 index 0000000..865027e Binary files /dev/null and b/img/DoF2_depth8_3000spp.png differ diff --git a/img/DoF_depth8_3000spp.png b/img/DoF_depth8_3000spp.png new file mode 100644 index 0000000..e7edf9d Binary files /dev/null and b/img/DoF_depth8_3000spp.png differ diff --git a/img/REFERENCE_cornell.5000samp.png b/img/REFERENCE_cornell.5000samp.png deleted file mode 100644 index 5ceb26e..0000000 Binary files a/img/REFERENCE_cornell.5000samp.png and /dev/null differ diff --git a/img/aa2_depth8_2000spp.png b/img/aa2_depth8_2000spp.png new file mode 100644 index 0000000..585910a Binary files /dev/null and b/img/aa2_depth8_2000spp.png differ diff --git a/img/aa3_depth8_2000spp.png b/img/aa3_depth8_2000spp.png new file mode 100644 index 0000000..c3469ba Binary files /dev/null and b/img/aa3_depth8_2000spp.png differ diff --git a/img/aa4_depth8_2000spp.png b/img/aa4_depth8_2000spp.png new file mode 100644 index 0000000..3c29af4 Binary files /dev/null and b/img/aa4_depth8_2000spp.png differ diff --git a/img/aa_depth8_2000spp.png b/img/aa_depth8_2000spp.png new file mode 100644 index 0000000..7441fa1 Binary files /dev/null and b/img/aa_depth8_2000spp.png differ diff --git a/img/cover_depth8_5000spp.png b/img/cover_depth8_5000spp.png new file mode 100644 index 0000000..049c827 Binary files /dev/null and b/img/cover_depth8_5000spp.png differ diff --git a/img/depth1.png b/img/depth1.png new file mode 100644 index 0000000..12b0149 Binary files /dev/null and b/img/depth1.png differ diff --git a/img/depth2.png b/img/depth2.png new file mode 100644 index 0000000..fb4eb01 Binary files /dev/null and b/img/depth2.png differ diff --git a/img/depth3.png b/img/depth3.png new file mode 100644 index 0000000..9964100 Binary files /dev/null and b/img/depth3.png differ diff --git a/img/depth4.png b/img/depth4.png new file mode 100644 index 0000000..ecdad8c Binary files /dev/null and b/img/depth4.png differ diff --git a/img/depth5.png b/img/depth5.png new file mode 100644 index 0000000..38f10ce Binary files /dev/null and b/img/depth5.png differ diff --git a/img/depth6.png b/img/depth6.png new file mode 100644 index 0000000..9bbde8a Binary files /dev/null and b/img/depth6.png differ diff --git a/img/depth7.png b/img/depth7.png new file mode 100644 index 0000000..8269c5b Binary files /dev/null and b/img/depth7.png differ diff --git a/img/depth8.png b/img/depth8.png new file mode 100644 index 0000000..030b338 Binary files /dev/null and b/img/depth8.png differ diff --git a/img/depth9.png b/img/depth9.png new file mode 100644 index 0000000..a17a2a6 Binary files /dev/null and b/img/depth9.png differ diff --git a/img/glossy2_depth8_3000spp.png b/img/glossy2_depth8_3000spp.png new file mode 100644 index 0000000..640ff4d Binary files /dev/null and b/img/glossy2_depth8_3000spp.png differ diff --git a/img/glossy_depth8_3000spp.png b/img/glossy_depth8_3000spp.png new file mode 100644 index 0000000..63be8b1 Binary files /dev/null and b/img/glossy_depth8_3000spp.png differ diff --git a/img/lambert_depth2_1000spp.png b/img/lambert_depth2_1000spp.png new file mode 100644 index 0000000..dbb8749 Binary files /dev/null and b/img/lambert_depth2_1000spp.png differ diff --git a/img/lambert_depth8_1000spp.png b/img/lambert_depth8_1000spp.png new file mode 100644 index 0000000..5a802fc Binary files /dev/null and b/img/lambert_depth8_1000spp.png differ diff --git a/img/mirror_depth2_1000spp.png b/img/mirror_depth2_1000spp.png new file mode 100644 index 0000000..f4ee500 Binary files /dev/null and b/img/mirror_depth2_1000spp.png differ diff --git a/img/mirror_depth8_1000spp.png b/img/mirror_depth8_1000spp.png new file mode 100644 index 0000000..5472b89 Binary files /dev/null and b/img/mirror_depth8_1000spp.png differ diff --git a/img/refraction2_depth8_1000spp.png b/img/refraction2_depth8_1000spp.png new file mode 100644 index 0000000..ff1949d Binary files /dev/null and b/img/refraction2_depth8_1000spp.png differ diff --git a/img/refraction_depth8_1000spp.png b/img/refraction_depth8_1000spp.png new file mode 100644 index 0000000..b4b7361 Binary files /dev/null and b/img/refraction_depth8_1000spp.png differ diff --git a/img/time_vs_depth.png b/img/time_vs_depth.png new file mode 100644 index 0000000..c849910 Binary files /dev/null and b/img/time_vs_depth.png differ diff --git a/scenes/cornell.txt b/scenes/cornell.txt index 83ff820..473e2f4 100644 --- a/scenes/cornell.txt +++ b/scenes/cornell.txt @@ -6,7 +6,7 @@ SPECRGB 0 0 0 REFL 0 REFR 0 REFRIOR 0 -EMITTANCE 5 +EMITTANCE 10 // Diffuse white MATERIAL 1 @@ -38,27 +38,79 @@ REFR 0 REFRIOR 0 EMITTANCE 0 -// Specular white +// Specular gold MATERIAL 4 -RGB .98 .98 .98 -SPECEX 0 +RGB .9 .7 .3 +SPECEX 100000 SPECRGB .98 .98 .98 REFL 1 REFR 0 REFRIOR 0 EMITTANCE 0 +// Refractive red +MATERIAL 5 +RGB .98 .28 .28 +SPECEX 5000 +SPECRGB .98 .98 .98 +REFL 0 +REFR 1 +REFRIOR 1.5 +EMITTANCE 0 + +// Lambert blue +MATERIAL 6 +RGB .28 .28 .98 +SPECEX 50000 +SPECRGB .98 .98 .98 +REFL 0 +REFR 0 +REFRIOR 1.5 +EMITTANCE 0 + +// Glossy gold +MATERIAL 7 +RGB .9 .7 .3 +SPECEX 1000 +SPECRGB .98 .98 .98 +REFL 1 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Refractive red +MATERIAL 8 +RGB .98 .28 .28 +SPECEX 5000 +SPECRGB .98 .98 .98 +REFL 0 +REFR 1 +REFRIOR 1.8 +EMITTANCE 0 + +// Lambert blue +MATERIAL 9 +RGB .28 .28 .98 +SPECEX 50000 +SPECRGB .98 .98 .98 +REFL 0 +REFR 0 +REFRIOR 1.5 +EMITTANCE 0 + // Camera CAMERA RES 800 800 FOVY 45 ITERATIONS 5000 -DEPTH 8 +DEPTH 9 FILE cornell -EYE 0.0 5 10.5 +EYE 0 5 10.5 LOOKAT 0 5 0 UP 0 1 0 - +LENSR 0.2 +FOCALD 7 +AA ON // Ceiling light OBJECT 0 @@ -108,10 +160,50 @@ TRANS 5 5 0 ROTAT 0 0 0 SCALE .01 10 10 -// Sphere +// Mirror Cube OBJECT 6 -sphere +cube material 4 -TRANS -1 4 -1 +TRANS -3 3 -2 +ROTAT 0 30 0 +SCALE 2 7 .5 + +// Transparent cube +OBJECT 7 +cube +material 5 +TRANS 0 3 -3 +ROTAT 0 0 0 +SCALE 2 7 .5 + +// Diffuse cube +OBJECT 8 +cube +material 6 +TRANS 3 3 -2 +ROTAT 0 -30 0 +SCALE 2 7 .5 + +// Mirror sphere +OBJECT 9 +sphere +material 7 +TRANS -3 1 2 +ROTAT 0 0 0 +SCALE 2 2 2 + +// Transparent sphere +OBJECT 10 +sphere +material 8 +TRANS 0 1 1 +ROTAT 0 0 0 +SCALE 2 2 2 + +// Diffuse sphere +OBJECT 11 +sphere +material 9 +TRANS 3 1 2 ROTAT 0 0 0 -SCALE 3 3 3 +SCALE 2 2 2 diff --git a/src/interactions.h b/src/interactions.h index 5ce3628..bc72dca 100644 --- a/src/interactions.h +++ b/src/interactions.h @@ -1,6 +1,7 @@ #pragma once #include "intersections.h" +#include // CHECKITOUT /** @@ -36,9 +37,9 @@ glm::vec3 calculateRandomDirectionInHemisphere( glm::vec3 perpendicularDirection2 = glm::normalize(glm::cross(normal, perpendicularDirection1)); - return up * normal + return glm::normalize(up * normal + cos(around) * over * perpendicularDirection1 - + sin(around) * over * perpendicularDirection2; + + sin(around) * over * perpendicularDirection2); } /** @@ -47,11 +48,11 @@ 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. This logic also applies to * combining other types of materias (such as refractive). - * + * * - Always take an even (50/50) split between a each effect (a diffuse bounce * and a specular bounce), but divide the resulting color of either branch * by its probability (0.5), to counteract the chance (0.5) of the branch @@ -67,13 +68,56 @@ glm::vec3 calculateRandomDirectionInHemisphere( * You may need to change the parameter list for your purposes! */ __host__ __device__ -void scatterRay( - PathSegment & pathSegment, - glm::vec3 intersect, - glm::vec3 normal, - const Material &m, - thrust::default_random_engine &rng) { +void scatterRay(PathSegment &path, const ShadeableIntersection &intersection, + const Material &material, thrust::default_random_engine &rng) { + // TODO: implement this. // A basic implementation of pure-diffuse shading will just call the // calculateRandomDirectionInHemisphere defined above. + + const Ray &ray = path.ray; + Ray newRay = path.ray; + thrust::uniform_real_distribution u01(0, 1); + + if (material.hasReflective > 0.f) { + // specular reflection + float theta = acos(pow(u01(rng), 1.f / (material.specular.exponent + 1.f))); + float phi = TWO_PI * u01(rng); + glm::vec3 mirror = glm::normalize(glm::reflect(ray.direction, + intersection.surfaceNormal)); + glm::vec3 up = {0.f, 0.f, 1.f}; + glm::vec3 axis = glm::normalize(glm::cross(up, mirror)); + float angle = acos(glm::dot(up, mirror)); + + newRay.direction = glm::vec3(cos(phi) * sin(theta), + sin(phi) * sin(theta), cos(theta)); + newRay.direction = glm::normalize( + glm::rotate(newRay.direction, angle, axis)); + } else if(material.hasRefractive > 0.f) { + // refraction + float cosine = glm::dot(intersection.surfaceNormal, ray.direction); + + if (glm::abs(cosine) < 1e-2f) { + newRay.direction = ray.direction; + } else if (cosine < 0.f) { + newRay.direction = glm::refract(ray.direction, + intersection.surfaceNormal, 1.f / material.indexOfRefraction); + } else { + newRay.direction = glm::refract(ray.direction, + -intersection.surfaceNormal, material.indexOfRefraction); + } + // total reflection + if (glm::length(newRay.direction) < 1e-3f){ + newRay.direction = glm::reflect(ray.direction, + -intersection.surfaceNormal); + } + } else { + // lambert reflection + newRay.direction = calculateRandomDirectionInHemisphere( + intersection.surfaceNormal, rng); + } + + newRay.origin = ray.origin + ray.direction * intersection.t + + newRay.direction * 1e-3f; + path.ray = newRay; } diff --git a/src/intersections.h b/src/intersections.h index 6f23872..98d4cf5 100644 --- a/src/intersections.h +++ b/src/intersections.h @@ -47,14 +47,16 @@ __host__ __device__ glm::vec3 multiplyMV(glm::mat4 m, glm::vec4 v) { */ __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.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)*/ { @@ -136,9 +138,9 @@ __host__ __device__ float sphereIntersectionTest(Geom sphere, Ray r, intersectionPoint = multiplyMV(sphere.transform, glm::vec4(objspaceIntersection, 1.f)); normal = glm::normalize(multiplyMV(sphere.invTranspose, glm::vec4(objspaceIntersection, 0.f))); - if (!outside) { - normal = -normal; - } + // if (!outside) { + // normal = -normal; + // } return glm::length(r.origin - intersectionPoint); } diff --git a/src/main.cpp b/src/main.cpp index fe8e85e..4eabd49 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -1,6 +1,8 @@ #include "main.h" #include "preview.h" #include +#include +using namespace std::chrono; static std::string startTimeString; @@ -22,6 +24,7 @@ glm::vec3 ogLookAt; // for recentering the camera Scene *scene; RenderState *renderState; int iteration; +int durationPerIteration; int width; int height; @@ -128,6 +131,9 @@ void runCuda() { } if (iteration < renderState->iterations) { + // tic + high_resolution_clock::time_point t1 = high_resolution_clock::now(); + uchar4 *pbo_dptr = NULL; iteration++; cudaGLMapBufferObject((void**)&pbo_dptr, pbo); @@ -138,6 +144,12 @@ void runCuda() { // unmap buffer object cudaGLUnmapBufferObject(pbo); + + // toc + cudaDeviceSynchronize(); + high_resolution_clock::time_point t2 = high_resolution_clock::now(); + duration d = duration_cast>(t2 - t1); + durationPerIteration = d.count(); } else { saveImage(); pathtraceFree(); diff --git a/src/main.h b/src/main.h index fdb7d5d..b8db37e 100644 --- a/src/main.h +++ b/src/main.h @@ -28,6 +28,7 @@ using namespace std; extern Scene* scene; extern int iteration; +extern int durationPerIteration; extern int width; extern int height; diff --git a/src/pathtrace.cu b/src/pathtrace.cu index c1ec122..7f1061f 100644 --- a/src/pathtrace.cu +++ b/src/pathtrace.cu @@ -15,9 +15,16 @@ #include "interactions.h" #define ERRORCHECK 1 - #define FILENAME (strrchr(__FILE__, '/') ? strrchr(__FILE__, '/') + 1 : __FILE__) #define checkCUDAError(msg) checkCUDAErrorFn(msg, FILENAME, __LINE__) + +struct AlivePathTest{ + __host__ __device__ + bool operator()(const PathSegment &path) { + return path.remainingBounces <= 0; + } +}; + void checkCUDAErrorFn(const char *msg, const char *file, int line) { #if ERRORCHECK cudaDeviceSynchronize(); @@ -111,172 +118,6 @@ void pathtraceFree() { checkCUDAError("pathtraceFree"); } -/** -* Generate PathSegments with rays from the camera through the screen into the -* scene, which is the first bounce of rays. -* -* Antialiasing - add rays for sub-pixel sampling -* motion blur - jitter rays "in time" -* lens effect - jitter ray origin positions based on a lens -*/ -__global__ void generateRayFromCamera(Camera cam, int iter, int traceDepth, PathSegment* pathSegments) -{ - 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); - PathSegment & segment = pathSegments[index]; - - segment.ray.origin = cam.position; - segment.color = glm::vec3(1.0f, 1.0f, 1.0f); - - // TODO: implement antialiasing by jittering the ray - segment.ray.direction = glm::normalize(cam.view - - cam.right * cam.pixelLength.x * ((float)x - (float)cam.resolution.x * 0.5f) - - cam.up * cam.pixelLength.y * ((float)y - (float)cam.resolution.y * 0.5f) - ); - - segment.pixelIndex = index; - segment.remainingBounces = traceDepth; - } -} - -// TODO: -// computeIntersections handles generating ray intersections ONLY. -// Generating new rays is handled in your shader(s). -// Feel free to modify the code below. -__global__ void computeIntersections( - int depth - , int num_paths - , PathSegment * pathSegments - , Geom * geoms - , int geoms_size - , ShadeableIntersection * intersections - ) -{ - int path_index = blockIdx.x * blockDim.x + threadIdx.x; - - if (path_index < num_paths) - { - PathSegment pathSegment = pathSegments[path_index]; - - float t; - glm::vec3 intersect_point; - glm::vec3 normal; - float t_min = FLT_MAX; - int hit_geom_index = -1; - bool outside = true; - - glm::vec3 tmp_intersect; - glm::vec3 tmp_normal; - - // naive parse through global geoms - - for (int i = 0; i < geoms_size; i++) - { - Geom & geom = geoms[i]; - - if (geom.type == CUBE) - { - t = boxIntersectionTest(geom, pathSegment.ray, tmp_intersect, tmp_normal, outside); - } - else if (geom.type == SPHERE) - { - t = sphereIntersectionTest(geom, pathSegment.ray, tmp_intersect, tmp_normal, outside); - } - // TODO: add more intersection tests here... triangle? metaball? CSG? - - // Compute the minimum t from the intersection tests to determine what - // scene geometry object was hit first. - if (t > 0.0f && t_min > t) - { - t_min = t; - hit_geom_index = i; - intersect_point = tmp_intersect; - normal = tmp_normal; - } - } - - if (hit_geom_index == -1) - { - intersections[path_index].t = -1.0f; - } - else - { - //The ray hits something - intersections[path_index].t = t_min; - intersections[path_index].materialId = geoms[hit_geom_index].materialid; - intersections[path_index].surfaceNormal = normal; - } - } -} - -// LOOK: "fake" shader demonstrating what you might do with the info in -// a ShadeableIntersection, as well as how to use thrust's random number -// generator. Observe that since the thrust random number generator basically -// adds "noise" to the iteration, the image should start off noisy and get -// cleaner as more iterations are computed. -// -// Note that this shader does NOT do a BSDF evaluation! -// Your shaders should handle that - this can allow techniques such as -// bump mapping. -__global__ void shadeFakeMaterial ( - int iter - , int num_paths - , ShadeableIntersection * shadeableIntersections - , PathSegment * pathSegments - , Material * materials - ) -{ - int idx = blockIdx.x * blockDim.x + threadIdx.x; - if (idx < num_paths) - { - ShadeableIntersection intersection = shadeableIntersections[idx]; - if (intersection.t > 0.0f) { // if the intersection exists... - // Set up the RNG - // LOOK: this is how you use thrust's RNG! Please look at - // makeSeededRandomEngine as well. - thrust::default_random_engine rng = makeSeededRandomEngine(iter, idx, 0); - thrust::uniform_real_distribution u01(0, 1); - - Material material = materials[intersection.materialId]; - glm::vec3 materialColor = material.color; - - // If the material indicates that the object was a light, "light" the ray - if (material.emittance > 0.0f) { - pathSegments[idx].color *= (materialColor * material.emittance); - } - // Otherwise, do some pseudo-lighting computation. This is actually more - // like what you would expect from shading in a rasterizer like OpenGL. - // TODO: replace this! you should be able to start with basically a one-liner - else { - float lightTerm = glm::dot(intersection.surfaceNormal, glm::vec3(0.0f, 1.0f, 0.0f)); - pathSegments[idx].color *= (materialColor * lightTerm) * 0.3f + ((1.0f - intersection.t * 0.02f) * materialColor) * 0.7f; - pathSegments[idx].color *= u01(rng); // apply some noise because why not - } - // If there was no intersection, color the ray black. - // Lots of renderers use 4 channel color, RGBA, where A = alpha, often - // used for opacity, in which case they can indicate "no opacity". - // This can be useful for post-processing and image compositing. - } else { - pathSegments[idx].color = glm::vec3(0.0f); - } - } -} - -// Add the current iteration's output to the overall image -__global__ void finalGather(int nPaths, glm::vec3 * image, PathSegment * iterationPaths) -{ - int index = (blockIdx.x * blockDim.x) + threadIdx.x; - - if (index < nPaths) - { - PathSegment iterationPath = iterationPaths[index]; - image[iterationPath.pixelIndex] += iterationPath.color; - } -} - /** * Wrapper for the __global__ call that sets up the kernel calls and does a ton * of memory management @@ -326,7 +167,8 @@ void pathtrace(uchar4 *pbo, int frame, int iter) { // TODO: perform one iteration of path tracing - generateRayFromCamera <<>>(cam, iter, traceDepth, dev_paths); + generateRayFromCamera<<>>( + cam, iter, traceDepth, dev_paths); checkCUDAError("generate camera ray"); int depth = 0; @@ -335,55 +177,62 @@ void pathtrace(uchar4 *pbo, int frame, int iter) { // --- PathSegment Tracing Stage --- // Shoot ray into scene, bounce between objects, push shading chunks - - bool iterationComplete = false; + bool iterationComplete = false; while (!iterationComplete) { - - // clean shading chunks - cudaMemset(dev_intersections, 0, pixelcount * sizeof(ShadeableIntersection)); - - // tracing - dim3 numblocksPathSegmentTracing = (num_paths + blockSize1d - 1) / blockSize1d; - computeIntersections <<>> ( - depth - , num_paths - , dev_paths - , dev_geoms - , hst_scene->geoms.size() - , dev_intersections - ); - checkCUDAError("trace one bounce"); - cudaDeviceSynchronize(); - depth++; - - - // TODO: - // --- Shading Stage --- - // Shade path segments based on intersections and generate new rays by - // evaluating the BSDF. - // Start off with just a big kernel that handles all the different - // materials you have in the scenefile. - // TODO: compare between directly shading the path segments and shading - // path segments that have been reshuffled to be contiguous in memory. - - shadeFakeMaterial<<>> ( - iter, - num_paths, - dev_intersections, - dev_paths, - dev_materials - ); - iterationComplete = true; // TODO: should be based off stream compaction results. + // clean shading chunks + cudaMemset(dev_intersections, 0, pixelcount * sizeof(ShadeableIntersection)); + + // tracing + dim3 numblocksPathSegmentTracing = (num_paths + blockSize1d - 1) / blockSize1d; + + computeIntersections<<>>( + depth, num_paths, dev_paths, dev_geoms, hst_scene->geoms.size(), + dev_intersections); + checkCUDAError("trace one bounce"); + cudaDeviceSynchronize(); + depth++; + + // TODO: + // --- Shading Stage --- + // Shade path segments based on intersections and generate new rays by + // evaluating the BSDF. + // Start off with just a big kernel that handles all the different + // materials you have in the scenefile. + // TODO: compare between directly shading the path segments and shading + // path segments that have been reshuffled to be contiguous in memory. + + // shadeFakeMaterial<<>>( + // iter, num_paths, dev_intersections, dev_paths, dev_materials); + + // working + shadeAllMaterial<<>>( + iter, num_paths, dev_intersections, dev_paths, dev_materials); + + gatherAndTerminate<<>>( + num_paths, dev_image, dev_paths, dev_intersections); + + scatterAliveRays<<>>( + iter, depth, num_paths, dev_paths, dev_intersections, + dev_materials); + + PathSegment *dev_path_end_new = thrust::remove_if(thrust::device, + dev_paths, dev_paths + num_paths, AlivePathTest()); + + if (depth >= traceDepth || dev_path_end_new == dev_paths) { + iterationComplete = true; // TODO: should be based off stream compaction results. + } else { + num_paths = dev_path_end_new - dev_paths; + } } - // Assemble this iteration and apply it to the image - dim3 numBlocksPixels = (pixelcount + blockSize1d - 1) / blockSize1d; - finalGather<<>>(num_paths, dev_image, dev_paths); + // Assemble this iteration and apply it to the image + dim3 numBlocksPixels = (pixelcount + blockSize1d - 1) / blockSize1d; /////////////////////////////////////////////////////////////////////////// // Send results to OpenGL buffer for rendering - sendImageToPBO<<>>(pbo, cam.resolution, iter, dev_image); + sendImageToPBO<<>>( + pbo, cam.resolution,iter, dev_image); // Retrieve image from GPU cudaMemcpy(hst_scene->state.image.data(), dev_image, @@ -391,3 +240,208 @@ void pathtrace(uchar4 *pbo, int frame, int iter) { checkCUDAError("pathtrace"); } + +/** +* Generate PathSegments with rays from the camera through the screen into the +* scene, which is the first bounce of rays. +* +* Antialiasing - add rays for sub-pixel sampling +* motion blur - jitter rays "in time" +* lens effect - jitter ray origin positions based on a lens +*/ +__global__ void generateRayFromCamera(Camera cam, int iter, int traceDepth, + PathSegment* pathSegments) { + + 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) { + return; + } + + int index = x + (y * cam.resolution.x); + float lensRadius = cam.lensRadius; + float focalDistance = cam.focalDistance; + thrust::default_random_engine rng = makeSeededRandomEngine( + iter, index, traceDepth); + thrust::uniform_real_distribution u01(0, 1); + PathSegment &segment = pathSegments[index]; + + // TODO: implement antialiasing by jittering the ray + float shiftX = 0.f; + float shiftY = 0.f; + + if (cam.isAA) { + shiftX = u01(rng) - .5f; + shiftY = u01(rng) - .5f; + } + + segment.ray.direction = glm::normalize(cam.view + - cam.right * cam.pixelLength.x * ((float)x - (float)cam.resolution.x * .5f + shiftX) + - cam.up * cam.pixelLength.y * ((float)y - (float)cam.resolution.y * .5f + shiftY)); + segment.color = glm::vec3(1.0f, 1.0f, 1.0f); + segment.pixelIndex = index; + segment.remainingBounces = traceDepth; + + // depth of field + if (lensRadius == 0.f) { + segment.ray.origin = cam.position; + } else { + float r = sqrtf(u01(rng)) * lensRadius; + float theta = u01(rng) * PI; + float ft = fabs(focalDistance / glm::dot(segment.ray.direction, cam.view)); + glm::vec3 pointFocus = cam.position + segment.ray.direction * ft; + + segment.ray.origin = cam.position + r * cosf(theta) * cam.right + + r * sinf(theta) * cam.up; + segment.ray.direction = glm::normalize(pointFocus - segment.ray.origin); + } +} + +// TODO: +// computeIntersections handles generating ray intersections ONLY. +// Generating new rays is handled in your shader(s). +// Feel free to modify the code below. +__global__ void computeIntersections(int depth, int num_paths, + PathSegment *pathSegments, Geom * geoms, int geoms_size, + ShadeableIntersection *intersections){ + + int path_index = blockIdx.x * blockDim.x + threadIdx.x; + + if (path_index >= num_paths) { + return; + } + + PathSegment pathSegment = pathSegments[path_index]; + + float t; + glm::vec3 intersect_point; + glm::vec3 normal; + float t_min = FLT_MAX; + int hit_geom_index = -1; + bool outside = true; + + glm::vec3 tmp_intersect; + glm::vec3 tmp_normal; + + // naive parse through global geoms + for (int i = 0; i < geoms_size; i++) { + + Geom &geom = geoms[i]; + + if (geom.type == CUBE) { + t = boxIntersectionTest(geom, pathSegment.ray, tmp_intersect, + tmp_normal, outside); + } else if (geom.type == SPHERE) { + t = sphereIntersectionTest(geom, pathSegment.ray, tmp_intersect, + tmp_normal, outside); + } + + // TODO: add more intersection tests here... triangle? metaball? CSG? + + // Compute the minimum t from the intersection tests to determine what + // scene geometry object was hit first. + if (t > 0.0f && t_min > t) { + t_min = t; + hit_geom_index = i; + intersect_point = tmp_intersect; + normal = tmp_normal; + } + } + + if (hit_geom_index == -1) { + intersections[path_index].t = -1.0f; + } else { + //The ray hits something + intersections[path_index].t = t_min; + intersections[path_index].materialId = geoms[hit_geom_index].materialid; + intersections[path_index].surfaceNormal = normal; + } +} + +__global__ void shadeAllMaterial(int iter, int num_paths, + ShadeableIntersection *shadeableIntersections, + PathSegment *pathSegments, Material *materials) { + + int idx = blockIdx.x * blockDim.x + threadIdx.x; + + if (idx >= num_paths) { + return; + } + + ShadeableIntersection intersection = shadeableIntersections[idx]; + + if (intersection.t > 0.0f) { + + thrust::default_random_engine rng = makeSeededRandomEngine(iter, idx, 0); + thrust::uniform_real_distribution u01(0, 1); + const Material &material = materials[intersection.materialId]; + + // If the material indicates that the object was a light, "light" the ray + if (material.emittance > 0.f) { + pathSegments[idx].remainingBounces = 0; + pathSegments[idx].color *= material.color * material.emittance; + } else { + // specular reflection + // lambert reflection + pathSegments[idx].color *= material.color; + } + } else { + // If there was no intersection, color the ray black. + // Lots of renderers use 4 channel color, RGBA, where A = alpha, often + // used for opacity, in which case they can indicate "no opacity". + // This can be useful for post-processing and image compositing. + pathSegments[idx].color = glm::vec3(0.0f); + } + + float colorRGB = pathSegments[idx].color.x + pathSegments[idx].color.y + + pathSegments[idx].color.z; + + if (colorRGB <= 0.f) { + pathSegments[idx].remainingBounces = 0; + } +} + +// Add the current iteration's output of terminated path to the overall image +__global__ void gatherAndTerminate(int num_paths, glm::vec3 * image, + PathSegment *dev_paths, ShadeableIntersection *dev_intersections) { + + int idx = (blockIdx.x * blockDim.x) + threadIdx.x; + + if (idx >= num_paths) { + return; + } + + PathSegment &path = dev_paths[idx]; + ShadeableIntersection &intersection = dev_intersections[idx]; + + path.remainingBounces -= 1; + if (path.remainingBounces <= 0) { + image[path.pixelIndex] += path.color; + } + if (intersection.t <= 1e-4f) { + path.remainingBounces = 0; + } +} + +__global__ void scatterAliveRays(int iter, int depth, int num_paths, + PathSegment *dev_paths, ShadeableIntersection *dev_intersections, + Material *dev_materials) { + + int idx = blockDim.x * blockIdx.x + threadIdx.x; + + if (idx >= num_paths) { + return; + } + + PathSegment &path = dev_paths[idx]; + ShadeableIntersection &intersection = dev_intersections[idx]; + Material &material = dev_materials[intersection.materialId]; + thrust::default_random_engine rng = makeSeededRandomEngine(iter, idx, depth); + + if (path.remainingBounces <= 0) { + return; + } + + scatterRay(path, intersection, material, rng); +} diff --git a/src/pathtrace.h b/src/pathtrace.h index 1241227..dcd0dc9 100644 --- a/src/pathtrace.h +++ b/src/pathtrace.h @@ -6,3 +6,21 @@ void pathtraceInit(Scene *scene); void pathtraceFree(); void pathtrace(uchar4 *pbo, int frame, int iteration); + +__global__ void gatherAndTerminate(int nPaths, glm::vec3 * image, + PathSegment *iterationPaths, ShadeableIntersection *intersctions); + +__global__ void shadeAllMaterial(int iter, int num_paths, + ShadeableIntersection *shadeableIntersections, + PathSegment *pathSegments, Material *materials); + +__global__ void computeIntersections(int depth, int num_paths, + PathSegment *pathSegments, Geom * geoms, int geoms_size, + ShadeableIntersection *intersections); + +__global__ void generateRayFromCamera(Camera cam, int iter, int traceDepth, + PathSegment* pathSegments); + +__global__ void scatterAliveRays(int iter, int depth, int num_paths, + PathSegment *dev_paths, ShadeableIntersection *dev_intersections, + Material *dev_materials); diff --git a/src/preview.cpp b/src/preview.cpp index 4eb0bc1..dd75473 100644 --- a/src/preview.cpp +++ b/src/preview.cpp @@ -14,7 +14,7 @@ std::string currentTimeString() { time_t now; time(&now); char buf[sizeof "0000-00-00_00-00-00z"]; - strftime(buf, sizeof buf, "%Y-%m-%d_%H-%M-%Sz", gmtime(&now)); + strftime(buf, sizeof buf, "%Y-%m-%d_%H-%M-%Sz", localtime(&now)); return std::string(buf); } @@ -173,7 +173,9 @@ void mainLoop() { glfwPollEvents(); runCuda(); - string title = "CIS565 Path Tracer | " + utilityCore::convertIntToString(iteration) + " Iterations"; + string title = "CIS565 Path Tracer | " + + utilityCore::convertIntToString(iteration) + " Iterations | " + + utilityCore::convertIntToString(durationPerIteration) + " ms"; glfwSetWindowTitle(window, title.c_str()); glBindBuffer(GL_PIXEL_UNPACK_BUFFER, pbo); diff --git a/src/scene.cpp b/src/scene.cpp index cbae043..2f8bce3 100644 --- a/src/scene.cpp +++ b/src/scene.cpp @@ -111,11 +111,12 @@ int Scene::loadCamera() { state.traceDepth = atoi(tokens[1].c_str()); } else if (strcmp(tokens[0].c_str(), "FILE") == 0) { state.imageName = tokens[1]; - } + } } string line; utilityCore::safeGetline(fp_in, line); + camera.isAA = false; while (!line.empty() && fp_in.good()) { vector tokens = utilityCore::tokenizeString(line); if (strcmp(tokens[0].c_str(), "EYE") == 0) { @@ -124,7 +125,13 @@ int Scene::loadCamera() { camera.lookAt = 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(), "LENSR") == 0) { + camera.lensRadius = atof(tokens[1].c_str()); + } else if (strcmp(tokens[0].c_str(), "FOCALD") == 0) { + camera.focalDistance = atof(tokens[1].c_str()); + } else if (strcmp(tokens[0].c_str(), "AA") == 0) { + camera.isAA = (tokens[1] == "ON"); + } utilityCore::safeGetline(fp_in, line); } diff --git a/src/sceneStructs.h b/src/sceneStructs.h index b38b820..9e10789 100644 --- a/src/sceneStructs.h +++ b/src/sceneStructs.h @@ -49,6 +49,9 @@ struct Camera { glm::vec3 right; glm::vec2 fov; glm::vec2 pixelLength; + float lensRadius; + float focalDistance; + bool isAA; }; struct RenderState { @@ -70,7 +73,7 @@ struct PathSegment { // 1) color contribution computation // 2) BSDF evaluation: generate a new ray struct ShadeableIntersection { - float t; - glm::vec3 surfaceNormal; - int materialId; + float t; + glm::vec3 surfaceNormal; + int materialId; };