diff --git a/README.md b/README.md index 110697c..2d44751 100644 --- a/README.md +++ b/README.md @@ -3,11 +3,64 @@ 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) +* Eric Chiu +* Tested on: Windows 10 Education, Intel(R) Xeon(R) CPU E5-1630 v4 @ 3.60GHz 32GB, NVIDIA GeForce GTX 1070 (SIGLAB) -### (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. +![](./img/procedural-shape-1.png) + +## Description + +This project implements a physically based pathtracer using CUDA programming and GPU hardware. Features include stream compaction, first bounce caching, material sorting, anti-aliasing, refraction with Frensel effects, and procedural shapes. + +## Stream Compaction + +Rendering without stream compaction was faster than rendering with stream compaction until the trace depth hit the 100s. This is probably because the time to perform stream compaction was more than the time it saved in the long run when the trace depth was under 100. When the trace depth was over 100, the time to perform stream compaction was marginal compared to the time it saved in the long run. + + +![](./img/stream-compaction.png) + + +## First Bounce Caching + +First bounce caching depends on the number of intersections found in the first bounce. As the image resolution increases, the more likely intersections are computed in the first bounce. We can see here that caching the first bounce slightly boosts performance, but not a lot. This is probably because the first bounce is only a small percentage of the number of bounces the pathtracer has to perform. + + +![](./img/first-bounce-caching.png) + + +## Material Sorting + +Sorting rays and path segments based on different materials surprisingly made the performance slower. The idea was to sort rays and path segments so that ones with the same material will be contiguous in memory. However, after further analysis, it seems like the time it takes to material sort outweights the amount of time it saves in later operations. + + +![](./img/material-sorting.png) + + +## Anti-Aliasing + +Anti-aliasing can be achieved by offsetting the rays by a small amount when it is first shot out of the screen. The pixel will then be an average of the offset rays, creating a smooth effect. We can see the effect take place in the two images below, on the edge on the sphere. + + +![](./img/aliased.png) + +![](./img/anti-aliased.png) + + +## Refraction + +We used Frensel effects with Schlick's approximation to simulate glass and refractive material. In the image below, we see spheres with diffuse, refractive, and reflective materials. Caustics, light focused in one area, is one result of refraction. + + +![](./img/refraction.png) + + +## Procedural Shapes + +Procedural shapes, specificially implicit surfaces, can be procedurally generated using signed distance functions. When a ray is shot, we can solve a distance function to determine whether the ray hit a implicit surface or not. If we do hit an implicit surface, we perform shading like for any other object. If we do not hit an implicit surface, we continue on with pathtracing. The two implicit surfaces implemented are shown below. + + +![](./img/procedural-shape-1.png) + +![](./img/procedural-shape-2.png) diff --git a/img/aliased.png b/img/aliased.png new file mode 100644 index 0000000..df61c71 Binary files /dev/null and b/img/aliased.png differ diff --git a/img/anti-aliased.png b/img/anti-aliased.png new file mode 100644 index 0000000..55409b1 Binary files /dev/null and b/img/anti-aliased.png differ diff --git a/img/cornell.2018-09-30_20-32-37z.5000samp.png b/img/cornell.2018-09-30_20-32-37z.5000samp.png new file mode 100644 index 0000000..c5370a7 Binary files /dev/null and b/img/cornell.2018-09-30_20-32-37z.5000samp.png differ diff --git a/img/cornell.2018-09-30_20-44-55z.5000samp.png b/img/cornell.2018-09-30_20-44-55z.5000samp.png new file mode 100644 index 0000000..ab5d9d1 Binary files /dev/null and b/img/cornell.2018-09-30_20-44-55z.5000samp.png differ diff --git a/img/first-bounce-caching.png b/img/first-bounce-caching.png new file mode 100644 index 0000000..ed8c4a4 Binary files /dev/null and b/img/first-bounce-caching.png differ diff --git a/img/material-sorting.png b/img/material-sorting.png new file mode 100644 index 0000000..c08c356 Binary files /dev/null and b/img/material-sorting.png differ diff --git a/img/procedural-shape-1.png b/img/procedural-shape-1.png new file mode 100644 index 0000000..53ce91a Binary files /dev/null and b/img/procedural-shape-1.png differ diff --git a/img/procedural-shape-2.png b/img/procedural-shape-2.png new file mode 100644 index 0000000..0a9a8ed Binary files /dev/null and b/img/procedural-shape-2.png differ diff --git a/img/refraction.png b/img/refraction.png new file mode 100644 index 0000000..b8d473e Binary files /dev/null and b/img/refraction.png differ diff --git a/img/stream-compaction.png b/img/stream-compaction.png new file mode 100644 index 0000000..7816fd7 Binary files /dev/null and b/img/stream-compaction.png differ diff --git a/scenes/cornell.txt b/scenes/cornell.txt index 83ff820..6e739d8 100644 --- a/scenes/cornell.txt +++ b/scenes/cornell.txt @@ -43,9 +43,9 @@ MATERIAL 4 RGB .98 .98 .98 SPECEX 0 SPECRGB .98 .98 .98 -REFL 1 +REFL 0 REFR 0 -REFRIOR 0 +REFRIOR 1.52 EMITTANCE 0 // Camera @@ -114,4 +114,4 @@ sphere material 4 TRANS -1 4 -1 ROTAT 0 0 0 -SCALE 3 3 3 +SCALE 1.5 1.5 1.5 diff --git a/scenes/cornell2.txt b/scenes/cornell2.txt new file mode 100644 index 0000000..614000e --- /dev/null +++ b/scenes/cornell2.txt @@ -0,0 +1,169 @@ +// Emissive material (light) +MATERIAL 0 +RGB 1 1 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 5 + +// Diffuse white +MATERIAL 1 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Diffuse red +MATERIAL 2 +RGB .85 .35 .35 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Diffuse green +MATERIAL 3 +RGB .35 .85 .35 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Glass +MATERIAL 4 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB .98 .98 .98 +REFL 1 +REFR 1 +REFRIOR 1.52 +EMITTANCE 0 + +// Mirror +MATERIAL 5 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB .98 .98 .98 +REFL 1 +REFR 0 +REFRIOR 1.52 +EMITTANCE 0 + +// Diffuse +MATERIAL 6 +RGB .35 .35 .98 +SPECEX 0 +SPECRGB .98 .98 .98 +REFL 0 +REFR 0 +REFRIOR 1.52 +EMITTANCE 0 + +// Camera +CAMERA +RES 800 800 +FOVY 45 +ITERATIONS 5000 +DEPTH 8 +FILE cornell +EYE 0.0 5 10.5 +LOOKAT 0 5 0 +UP 0 1 0 + + +// Ceiling light +OBJECT 0 +cube +material 0 +TRANS 0 10 0 +ROTAT 0 0 0 +SCALE 3 .3 3 + +// Floor +OBJECT 1 +cube +material 1 +TRANS 0 0 0 +ROTAT 0 0 0 +SCALE 10 .01 10 + +// Ceiling +OBJECT 2 +cube +material 1 +TRANS 0 10 0 +ROTAT 0 0 90 +SCALE .01 10 10 + +// Back wall +OBJECT 3 +cube +material 1 +TRANS 0 5 -5 +ROTAT 0 90 0 +SCALE .01 10 10 + +// Left wall +OBJECT 4 +cube +material 2 +TRANS -5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +// Right wall +OBJECT 5 +cube +material 3 +TRANS 5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +// Sphere +OBJECT 6 +sphere +material 4 +TRANS -3 7 -1 +ROTAT 0 0 0 +SCALE 2.5 2.5 2.5 + +// Sphere +OBJECT 7 +sphere +material 4 +TRANS 0.5 1.75 2 +ROTAT 0 0 0 +SCALE 0.75 0.75 0.75 + +// Sphere +OBJECT 8 +sphere +material 5 +TRANS -2 3 -1 +ROTAT 0 0 0 +SCALE 4 4 4 + +// Sphere +OBJECT 9 +sphere +material 4 +TRANS 2.5 6 -1 +ROTAT 0 0 0 +SCALE 3.5 3.5 3.5 + +// Sphere +OBJECT 10 +sphere +material 6 +TRANS 2.5 2 -1 +ROTAT 0 0 0 +SCALE 1.5 1.5 1.5 diff --git a/scenes/csg1.txt b/scenes/csg1.txt new file mode 100644 index 0000000..f1224e0 --- /dev/null +++ b/scenes/csg1.txt @@ -0,0 +1,117 @@ +// Emissive material (light) +MATERIAL 0 +RGB 1 1 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 5 + +// Diffuse white +MATERIAL 1 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Diffuse red +MATERIAL 2 +RGB .85 .35 .35 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Diffuse green +MATERIAL 3 +RGB .35 .85 .35 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Specular white +MATERIAL 4 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB .98 .98 .98 +REFL 1 +REFR 1 +REFRIOR 1.52 +EMITTANCE 0 + +// Camera +CAMERA +RES 800 800 +FOVY 45 +ITERATIONS 5000 +DEPTH 8 +FILE cornell +EYE 0.0 5 10.5 +LOOKAT 0 5 0 +UP 0 1 0 + + +// Ceiling light +OBJECT 0 +cube +material 0 +TRANS 0 10 0 +ROTAT 0 0 0 +SCALE 3 .3 3 + +// Floor +OBJECT 1 +cube +material 1 +TRANS 0 0 0 +ROTAT 0 0 0 +SCALE 10 .01 10 + +// Ceiling +OBJECT 2 +cube +material 1 +TRANS 0 10 0 +ROTAT 0 0 90 +SCALE .01 10 10 + +// Back wall +OBJECT 3 +cube +material 1 +TRANS 0 5 -5 +ROTAT 0 90 0 +SCALE .01 10 10 + +// Left wall +OBJECT 4 +cube +material 2 +TRANS -5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +// Right wall +OBJECT 5 +cube +material 3 +TRANS 5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +// CSG1 +OBJECT 6 +csg1 +material 4 +TRANS 0 3 0 +ROTAT 0 0 0 +SCALE 1 1 1 diff --git a/scenes/csg2.txt b/scenes/csg2.txt new file mode 100644 index 0000000..b3c786c --- /dev/null +++ b/scenes/csg2.txt @@ -0,0 +1,117 @@ +// Emissive material (light) +MATERIAL 0 +RGB 1 1 1 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 5 + +// Diffuse white +MATERIAL 1 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Diffuse red +MATERIAL 2 +RGB .85 .35 .35 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Diffuse green +MATERIAL 3 +RGB .35 .85 .35 +SPECEX 0 +SPECRGB 0 0 0 +REFL 0 +REFR 0 +REFRIOR 0 +EMITTANCE 0 + +// Specular white +MATERIAL 4 +RGB .98 .98 .98 +SPECEX 0 +SPECRGB .98 .98 .98 +REFL 1 +REFR 1 +REFRIOR 1.52 +EMITTANCE 0 + +// Camera +CAMERA +RES 800 800 +FOVY 45 +ITERATIONS 5000 +DEPTH 8 +FILE cornell +EYE 0.0 5 10.5 +LOOKAT 0 5 0 +UP 0 1 0 + + +// Ceiling light +OBJECT 0 +cube +material 0 +TRANS 0 10 0 +ROTAT 0 0 0 +SCALE 3 .3 3 + +// Floor +OBJECT 1 +cube +material 1 +TRANS 0 0 0 +ROTAT 0 0 0 +SCALE 10 .01 10 + +// Ceiling +OBJECT 2 +cube +material 1 +TRANS 0 10 0 +ROTAT 0 0 90 +SCALE .01 10 10 + +// Back wall +OBJECT 3 +cube +material 1 +TRANS 0 5 -5 +ROTAT 0 90 0 +SCALE .01 10 10 + +// Left wall +OBJECT 4 +cube +material 2 +TRANS -5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +// Right wall +OBJECT 5 +cube +material 3 +TRANS 5 5 0 +ROTAT 0 0 0 +SCALE .01 10 10 + +// CSG2 +OBJECT 6 +csg2 +material 4 +TRANS 0.20 3.75 3 +ROTAT 30 -70 90 +SCALE 0.45 0.45 0.45 diff --git a/src/interactions.h b/src/interactions.h index 5ce3628..0811969 100644 --- a/src/interactions.h +++ b/src/interactions.h @@ -41,6 +41,60 @@ glm::vec3 calculateRandomDirectionInHemisphere( + sin(around) * over * perpendicularDirection2; } +__forceinline__ +__host__ __device__ +void reflective( + PathSegment &pathSegment, + glm::vec3 intersect, + glm::vec3 normal, + const Material &m, + thrust::default_random_engine &rng) +{ + glm::vec3 dir = glm::normalize(pathSegment.ray.direction); + glm::vec3 nor = glm::normalize(normal); + pathSegment.ray.direction = glm::reflect(dir, nor); + pathSegment.color *= m.color; + pathSegment.ray.origin = intersect + (.001f) * pathSegment.ray.direction; +} + +__forceinline__ +__host__ __device__ +void refractive( + PathSegment &pathSegment, + glm::vec3 intersect, + glm::vec3 normal, + const Material &m, + thrust::default_random_engine &rng) +{ + glm::vec3 dir = glm::normalize(pathSegment.ray.direction); + glm::vec3 nor = glm::normalize(normal); + float ior = m.indexOfRefraction; + + // if ray is in the same as normal direction, then ray is inside object + // surface normal must be flipped so that surface is facing ray + if (glm::dot(dir, nor) > 0) + { + nor = -nor; + } + else + { + ior = 1.0f / ior; + } + + // check for total internal reflection + + if (glm::length(pathSegment.ray.direction) < 0.01f) { + pathSegment.ray.direction = glm::reflect(dir, nor); + } + else + { + pathSegment.ray.direction = glm::refract(dir, nor, ior); + } + + pathSegment.color *= m.color; + pathSegment.ray.origin = intersect + (.001f) * pathSegment.ray.direction; +} + /** * Scatter a ray with some probabilities according to the material properties. * For example, a diffuse surface scatters in a cosine-weighted hemisphere. @@ -68,12 +122,82 @@ glm::vec3 calculateRandomDirectionInHemisphere( */ __host__ __device__ void scatterRay( - PathSegment & pathSegment, + PathSegment &pathSegment, glm::vec3 intersect, glm::vec3 normal, const Material &m, - thrust::default_random_engine &rng) { + thrust::default_random_engine &rng) +{ // TODO: implement this. // A basic implementation of pure-diffuse shading will just call the // calculateRandomDirectionInHemisphere defined above. + + glm::vec3 dir = glm::normalize(pathSegment.ray.direction); + glm::vec3 nor = normal; + + if (m.hasReflective && m.hasRefractive) + { + thrust::uniform_real_distribution u01(0, 1); + + float cosThetaI = glm::clamp(glm::dot(dir, nor), -1.0f, 1.0f); + float etaI = 1.f; + float etaT = m.indexOfRefraction; + + // outside object + if (cosThetaI < 0.f) { + cosThetaI = -cosThetaI; + pathSegment.ray.origin = intersect + (.001f) * dir; + } + // inside object + else + { + float temp = etaI; + etaI = etaT; + etaT = temp; + nor = -nor; + pathSegment.ray.origin = intersect + (.001f) * dir; + } + + float sinThetaI = std::sqrt(std::max(0.0f, 1.0f - cosThetaI * cosThetaI)); + float sinThetaT = etaI / etaT * sinThetaI; + + // total internal reflection + if (sinThetaT >= 1.0f) { + pathSegment.ray.direction = glm::reflect(pathSegment.ray.direction, nor); + pathSegment.color *= m.specular.color; + } + // use fresnel schlick's approximation + else { + + // Schlick's Approx + float r0 = glm::pow((etaI - etaT) / (etaI + etaT), 2.f); + float fresnel = r0 + (1.0f - r0) * pow(1.f - glm::abs(glm::dot(normal, pathSegment.ray.direction)), 5.f); + + if (fresnel < u01(rng)) { + refractive(pathSegment, intersect, normal, m, rng); + } + else { + reflective(pathSegment, intersect, normal, m, rng); + } + } + + } + // reflective case + else if (m.hasReflective) + { + reflective(pathSegment, intersect, normal, m, rng); + } + // refractive case + else if (m.hasRefractive) + { + refractive(pathSegment, intersect, normal, m, rng); + } + // diffuse case + else + { + pathSegment.ray.direction = calculateRandomDirectionInHemisphere(nor, rng); + pathSegment.color *= m.color; + pathSegment.ray.origin = intersect + (.001f) * pathSegment.ray.direction; + } + } diff --git a/src/intersections.h b/src/intersections.h index 6f23872..028004b 100644 --- a/src/intersections.h +++ b/src/intersections.h @@ -142,3 +142,158 @@ __host__ __device__ float sphereIntersectionTest(Geom sphere, Ray r, return glm::length(r.origin - intersectionPoint); } + + + +__host__ __device__ float csg1SDF(glm::vec3 p) +{ + float x4 = p.x * p.x * p.x * p.x; + float x2 = p.x * p.x; + float y4 = p.y * p.y * p.y * p.y; + float y2 = p.y * p.y; + float z4 = p.z * p.z * p.z * p.z; + float z2 = p.z * p.z; + return x4 - 5 * x2 + y4 - 5 * y2 + z4 - 5 * z2 + 11.8; +} + +__host__ __device__ float csg1Raytrace(glm::vec3 cam, glm::vec3 ray, float maxdist, bool &outside) +{ + float BIGSTEPSIZE = 0.1; + float SMALLSTEPSIZE = 0.02; + float t = 0.0f; + float step = BIGSTEPSIZE; + + for (int i = 0; i < 700; i++) { + glm::vec3 p = cam + ray * t; + float distance = csg1SDF(p); + + if (distance < 0.001) { + + //step back once, and increment slowly + t -= step; + step = SMALLSTEPSIZE; + + for (int i = 0; i < 10; i++) + { + p = cam + ray * t; + distance = csg1SDF(p); + if (distance < 0.001) { + t -= step; + p = cam + ray * t; + // InitializeIntersection(isect, t, p); + return t; + } + t += step; + } + return 0; + } + t += step; + } + return 0; +} + +__host__ __device__ glm::vec3 getCsg1Normal(glm::vec3 p) +{ + return glm::normalize(glm::vec3( + csg1SDF(glm::vec3(p.x + EPSILON, p.y, p.z)) - csg1SDF(glm::vec3(p.x - EPSILON, p.y, p.z)), + csg1SDF(glm::vec3(p.x, p.y + EPSILON, p.z)) - csg1SDF(glm::vec3(p.x, p.y - EPSILON, p.z)), + csg1SDF(glm::vec3(p.x, p.y, p.z + EPSILON)) - csg1SDF(glm::vec3(p.x, p.y, p.z - EPSILON)) + )); +} + +__host__ __device__ float csg1IntersectionTest(Geom surface, Ray r, + glm::vec3 &intersectionPoint, glm::vec3 &normal, bool &outside) +{ + glm::vec3 pt = multiplyMV(surface.inverseTransform, glm::vec4(r.origin, 1.0f)); + glm::vec3 dir = multiplyMV(surface.inverseTransform, glm::vec4(r.direction, 0.0f)); + + // raytrace to get t value + float t = csg1Raytrace(pt, dir, 100.f, outside); + + // calculate point on surface using ray and t value + glm::vec3 objPt = pt + t * dir; + intersectionPoint = multiplyMV(surface.transform, glm::vec4(objPt, 1.0f)); + + // calculate normal using gradient + normal = glm::normalize(multiplyMV(surface.invTranspose, glm::vec4(getCsg1Normal(objPt), 0.0f))); + if (!outside) normal = -normal; + return t; +} + + + +__host__ __device__ float csg2SDF(glm::vec3 p) +{ + float k = 5.0; + float a = 0.95; + float b = 0.5; + + float x2 = p.x * p.x; + float y2 = p.y * p.y; + float z2 = p.z * p.z; + return (x2 + y2 + z2 - a*k*k) * (x2 + y2 + z2 - a*k*k) - b * ((p.z - k)*(p.z - k) - 2*p.x*p.x) * ((p.z + k)*(p.z + k) - 2 *p.y*p.y); +} + +__host__ __device__ float csg2Raytrace(glm::vec3 cam, glm::vec3 ray, float maxdist, bool &outside) +{ + float BIGSTEPSIZE = 0.1; + float SMALLSTEPSIZE = 0.02; + float t = 0.0f; + float step = BIGSTEPSIZE; + + for (int i = 0; i < 700; i++) { + glm::vec3 p = cam + ray * t; + float distance = csg2SDF(p); + + if (distance < 0.001) { + + //step back once, and increment slowly + t -= step; + step = SMALLSTEPSIZE; + + for (int i = 0; i < 10; i++) + { + p = cam + ray * t; + distance = csg2SDF(p); + if (distance < 0.001) { + t -= step; + p = cam + ray * t; + // InitializeIntersection(isect, t, p); + return t; + } + t += step; + } + return 0; + } + t += step; + } + return 0; +} + +__host__ __device__ glm::vec3 getCsg2Normal(glm::vec3 p) +{ + return glm::normalize(glm::vec3( + csg2SDF(glm::vec3(p.x + EPSILON, p.y, p.z)) - csg2SDF(glm::vec3(p.x - EPSILON, p.y, p.z)), + csg2SDF(glm::vec3(p.x, p.y + EPSILON, p.z)) - csg2SDF(glm::vec3(p.x, p.y - EPSILON, p.z)), + csg2SDF(glm::vec3(p.x, p.y, p.z + EPSILON)) - csg2SDF(glm::vec3(p.x, p.y, p.z - EPSILON)) + )); +} + +__host__ __device__ float csg2IntersectionTest(Geom surface, Ray r, + glm::vec3 &intersectionPoint, glm::vec3 &normal, bool &outside) +{ + glm::vec3 pt = multiplyMV(surface.inverseTransform, glm::vec4(r.origin, 1.0f)); + glm::vec3 dir = multiplyMV(surface.inverseTransform, glm::vec4(r.direction, 0.0f)); + + // raytrace to get t value + float t = csg2Raytrace(pt, dir, 100.f, outside); + + // calculate point on surface using ray and t value + glm::vec3 objPt = pt + t * dir; + intersectionPoint = multiplyMV(surface.transform, glm::vec4(objPt, 1.0f)); + + // calculate normal using gradient + normal = glm::normalize(multiplyMV(surface.invTranspose, glm::vec4(getCsg2Normal(objPt), 0.0f))); + if (!outside) normal = -normal; + return t; +} \ No newline at end of file diff --git a/src/pathtrace.cu b/src/pathtrace.cu index c1ec122..ba586a4 100644 --- a/src/pathtrace.cu +++ b/src/pathtrace.cu @@ -4,6 +4,7 @@ #include #include #include +#include #include "sceneStructs.h" #include "scene.h" @@ -46,7 +47,7 @@ thrust::default_random_engine makeSeededRandomEngine(int iter, int index, int de //Kernel that writes the image to the OpenGL PBO directly. __global__ void sendImageToPBO(uchar4* pbo, glm::ivec2 resolution, - int iter, glm::vec3* image) { + int iter, glm::vec3* image) { int x = (blockIdx.x * blockDim.x) + threadIdx.x; int y = (blockIdx.y * blockDim.y) + threadIdx.y; @@ -55,9 +56,9 @@ __global__ void sendImageToPBO(uchar4* pbo, glm::ivec2 resolution, 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); + 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; @@ -75,6 +76,8 @@ static PathSegment * dev_paths = NULL; static ShadeableIntersection * dev_intersections = NULL; // TODO: static variables for device memory, any extra info you need, etc // ... +// TODO: Part 1 - Caching first bounce intersections +static ShadeableIntersection * dev_first_intersections = NULL; void pathtraceInit(Scene *scene) { hst_scene = scene; @@ -84,29 +87,38 @@ void pathtraceInit(Scene *scene) { cudaMalloc(&dev_image, pixelcount * sizeof(glm::vec3)); cudaMemset(dev_image, 0, pixelcount * sizeof(glm::vec3)); - cudaMalloc(&dev_paths, pixelcount * sizeof(PathSegment)); + cudaMalloc(&dev_paths, pixelcount * sizeof(PathSegment)); - cudaMalloc(&dev_geoms, scene->geoms.size() * sizeof(Geom)); - cudaMemcpy(dev_geoms, scene->geoms.data(), scene->geoms.size() * sizeof(Geom), cudaMemcpyHostToDevice); + cudaMalloc(&dev_geoms, scene->geoms.size() * sizeof(Geom)); + cudaMemcpy(dev_geoms, scene->geoms.data(), scene->geoms.size() * sizeof(Geom), cudaMemcpyHostToDevice); - cudaMalloc(&dev_materials, scene->materials.size() * sizeof(Material)); - cudaMemcpy(dev_materials, scene->materials.data(), scene->materials.size() * sizeof(Material), cudaMemcpyHostToDevice); + cudaMalloc(&dev_materials, scene->materials.size() * sizeof(Material)); + cudaMemcpy(dev_materials, scene->materials.data(), scene->materials.size() * sizeof(Material), cudaMemcpyHostToDevice); - cudaMalloc(&dev_intersections, pixelcount * sizeof(ShadeableIntersection)); - cudaMemset(dev_intersections, 0, pixelcount * sizeof(ShadeableIntersection)); + cudaMalloc(&dev_intersections, pixelcount * sizeof(ShadeableIntersection)); + cudaMemset(dev_intersections, 0, pixelcount * sizeof(ShadeableIntersection)); // TODO: initialize any extra device memeory you need + // TODO: Part 1 - Caching first bounce intersections + #if CACHING + cudaMalloc(&dev_first_intersections, pixelcount * sizeof(ShadeableIntersection)); + #endif checkCUDAError("pathtraceInit"); } void pathtraceFree() { cudaFree(dev_image); // no-op if dev_image is null - cudaFree(dev_paths); - cudaFree(dev_geoms); - cudaFree(dev_materials); - cudaFree(dev_intersections); - // TODO: clean up any extra device memory you created + cudaFree(dev_paths); + cudaFree(dev_geoms); + cudaFree(dev_materials); + cudaFree(dev_intersections); + + // clean up any extra device memory you created + // TODO: Part 1 - Cache first bounce intersections + #if CACHING + cudaFree(dev_first_intersections); + #endif checkCUDAError("pathtraceFree"); } @@ -121,25 +133,36 @@ void pathtraceFree() { */ __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) - ); + int x = (blockIdx.x * blockDim.x) + threadIdx.x; + int y = (blockIdx.y * blockDim.y) + threadIdx.y; - segment.pixelIndex = index; - segment.remainingBounces = traceDepth; - } + 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: Part 2 - implement antialiasing by jittering the ray + float xOffset = 0.0f; + float yOffset = 0.0f; + #if !CACHING + // use random number generator to add offset to x and y + thrust::default_random_engine rng = makeSeededRandomEngine(iter, index, traceDepth); + thrust::uniform_real_distribution u01(0, 1); + xOffset = u01(rng); + yOffset = u01(rng); + #endif + + segment.ray.direction = glm::normalize(cam.view + - cam.right * cam.pixelLength.x * ((float)(x + xOffset) - (float)cam.resolution.x * 0.5f) + - cam.up * cam.pixelLength.y * ((float)(y + yOffset) - (float)cam.resolution.y * 0.5f) + ); + + segment.pixelIndex = index; + segment.remainingBounces = traceDepth; + } } // TODO: @@ -147,69 +170,77 @@ __global__ void generateRayFromCamera(Camera cam, int iter, int traceDepth, Path // 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 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; - } - } + 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? + else if (geom.type == CSG1) + { + t = csg1IntersectionTest(geom, pathSegment.ray, tmp_intersect, tmp_normal, outside); + } + else if (geom.type == CSG2) + { + t = csg2IntersectionTest(geom, pathSegment.ray, tmp_intersect, tmp_normal, outside); + } + + // 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 @@ -221,62 +252,99 @@ __global__ void computeIntersections( // 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 - ) +__global__ void shadeFakeMaterial( + int iter + , int num_paths + , ShadeableIntersection * shadeableIntersections + , PathSegment * pathSegments + , Material * materials + , int depth +) { - 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); + int idx = blockIdx.x * blockDim.x + threadIdx.x; + + // not needed if using compact, needed if not using compact + // without compact, you don't know which rays are finished + #if !COMPACT + if (pathSegments[idx].remainingBounces == 0) return; + #endif + + 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, depth); + 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); + pathSegments[idx].remainingBounces = 0; + } + // 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 + // TODO: Part 1 - Shading kernel with BSDF evaluation + glm::vec3 intersectionPoint = getPointOnRay(pathSegments[idx].ray, intersection.t); + scatterRay(pathSegments[idx], intersectionPoint, intersection.surfaceNormal, material, rng); + pathSegments[idx].remainingBounces--; + } + // 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); + pathSegments[idx].remainingBounces = 0; + } } - } } // 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; + int index = (blockIdx.x * blockDim.x) + threadIdx.x; - if (index < nPaths) - { - PathSegment iterationPath = iterationPaths[index]; - image[iterationPath.pixelIndex] += iterationPath.color; - } + if (index < nPaths) + { + PathSegment iterationPath = iterationPaths[index]; + image[iterationPath.pixelIndex] += iterationPath.color; + } } +// TODO: Part 1 - Stream Compaction +// A predicate used for thrust::partition function to tell how to partition +struct streamCompactPredicate { + + streamCompactPredicate() {}; + + __host__ __device__ bool operator()(const PathSegment &s) { + return (s.remainingBounces > 0); + } +}; + +// TODO: Part 1 - Sorting rays, pathSegments, intersections +// A predicate used for thrust::partition function to tell how to partition +struct sortPredicate { + + sortPredicate() {}; + + __host__ __device__ bool operator()(const ShadeableIntersection &i1, const ShadeableIntersection &i2) { + return (i1.materialId < i2.materialId); + } +}; + /** * Wrapper for the __global__ call that sets up the kernel calls and does a ton * of memory management @@ -286,14 +354,14 @@ void pathtrace(uchar4 *pbo, int frame, int iter) { const Camera &cam = hst_scene->state.camera; const int pixelcount = cam.resolution.x * cam.resolution.y; - // 2D block for generating ray from camera + // 2D block for generating ray from camera const dim3 blockSize2d(8, 8); const dim3 blocksPerGrid2d( - (cam.resolution.x + blockSize2d.x - 1) / blockSize2d.x, - (cam.resolution.y + blockSize2d.y - 1) / blockSize2d.y); + (cam.resolution.x + blockSize2d.x - 1) / blockSize2d.x, + (cam.resolution.y + blockSize2d.y - 1) / blockSize2d.y); - // 1D block for path tracing - const int blockSize1d = 128; + // 1D block for path tracing + const int blockSize1d = 128; /////////////////////////////////////////////////////////////////////////// @@ -326,68 +394,139 @@ void pathtrace(uchar4 *pbo, int frame, int iter) { // TODO: perform one iteration of path tracing - generateRayFromCamera <<>>(cam, iter, traceDepth, dev_paths); - checkCUDAError("generate camera ray"); - - int depth = 0; - PathSegment* dev_path_end = dev_paths + pixelcount; - int num_paths = dev_path_end - dev_paths; - - // --- PathSegment Tracing Stage --- - // Shoot ray into scene, bounce between objects, push shading chunks - - 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. - } - - // Assemble this iteration and apply it to the image - dim3 numBlocksPixels = (pixelcount + blockSize1d - 1) / blockSize1d; - finalGather<<>>(num_paths, dev_image, dev_paths); + generateRayFromCamera << > > (cam, iter, traceDepth, dev_paths); + checkCUDAError("generate camera ray"); + + int depth = 0; + PathSegment* dev_path_end = dev_paths + pixelcount; + int num_paths = dev_path_end - dev_paths; + + #if TIMING + using time_point_t = std::chrono::high_resolution_clock::time_point; + time_point_t startTime = std::chrono::high_resolution_clock::now(); + #endif + + // --- PathSegment Tracing Stage --- + // Shoot ray into scene, bounce between objects, push shading chunks + + bool iterationComplete = false; + while (!iterationComplete) { + + // clean shading chunks + cudaMemset(dev_intersections, 0, pixelcount * sizeof(ShadeableIntersection)); + + dim3 numblocksPathSegmentTracing = (num_paths + blockSize1d - 1) / blockSize1d; + + #if CACHING + if ((iter == 1 && depth == 0) || depth > 0) + { + computeIntersections << > > ( + depth, + num_paths, + dev_paths, + dev_geoms, + hst_scene->geoms.size(), + dev_intersections + ); + checkCUDAError("trace one bounce"); + } + // TODO: Part 1 - Caching first bounce intersections + // iteration is the sample number of the pixel + // iteration == 1 means the first ray shot out of pixel + // depth is the the bounce number + if (iter == 1 && depth == 0) + { + // store computed intersection into cache + cudaMemcpy(dev_first_intersections, dev_intersections, pixelcount * sizeof(ShadeableIntersection), cudaMemcpyDeviceToDevice); + } + // TODO: Part 1 - Caching first bounce intersection + // if ray is not the first ray shot out of pixel, and + // is the first bounce, use the saved cache + else if (iter > 1 && depth == 0) + { + cudaMemcpy(dev_intersections, dev_first_intersections, pixelcount * sizeof(ShadeableIntersection), cudaMemcpyDeviceToDevice); + } + cudaDeviceSynchronize(); + depth++; + // NO CACHING + #else + // tracing + computeIntersections<<>> ( + depth, + num_paths, + dev_paths, + dev_geoms, + hst_scene->geoms.size(), + dev_intersections + ); + checkCUDAError("trace one bounce"); + cudaDeviceSynchronize(); + depth++; + #endif + + // TODO: Part 1 - Sorting rays, pathSegments, intersections + #if SORTING + #if SORTTIMING + using time_point_s = std::chrono::high_resolution_clock::time_point; + time_point_s startTime2 = std::chrono::high_resolution_clock::now(); + #endif + thrust::sort_by_key(thrust::device, dev_intersections, dev_intersections + num_paths, dev_paths, sortPredicate()); + #if SORTTIMING + cudaDeviceSynchronize(); + time_point_s endTime2 = std::chrono::high_resolution_clock::now(); + std::chrono::duration dur2 = endTime2 - startTime2; + float elapsedTime2 = static_cast(dur2.count()); + std::cout << "sorting time: " << elapsedTime2 << " milliseconds" << std::endl; + #endif + #endif + + // 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, + depth + ); + + #if COMPACT + // TODO: Part 1 - Stream Compaction + PathSegment* endSegment = thrust::partition(thrust::device, dev_paths, dev_paths + num_paths, streamCompactPredicate()); + num_paths = endSegment - dev_paths; + iterationComplete = (depth >= traceDepth) || num_paths <= 0; // TODO: iterationComplete should be based off stream compaction results. + #else + iterationComplete = (depth >= traceDepth); + #endif + } + + #if TIMING + cudaDeviceSynchronize(); + time_point_t endTime = std::chrono::high_resolution_clock::now(); + std::chrono::duration dur = endTime - startTime; + float elapsedTime = static_cast(dur.count()); + std::cout << "elapsed time: " << elapsedTime << " milliseconds" << std::endl; + #endif + + // Assemble this iteration and apply it to the image + dim3 numBlocksPixels = (pixelcount + blockSize1d - 1) / blockSize1d; + finalGather << > > (pixelcount, dev_image, dev_paths); /////////////////////////////////////////////////////////////////////////// // 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, - pixelcount * sizeof(glm::vec3), cudaMemcpyDeviceToHost); + pixelcount * sizeof(glm::vec3), cudaMemcpyDeviceToHost); checkCUDAError("pathtrace"); } diff --git a/src/pathtrace.h b/src/pathtrace.h index 1241227..2a07d93 100644 --- a/src/pathtrace.h +++ b/src/pathtrace.h @@ -3,6 +3,12 @@ #include #include "scene.h" +#define COMPACT 0 +#define SORTING 0 +#define CACHING 0 +#define TIMING 0 +#define SORTTIMING 0 + void pathtraceInit(Scene *scene); void pathtraceFree(); void pathtrace(uchar4 *pbo, int frame, int iteration); diff --git a/src/scene.cpp b/src/scene.cpp index cbae043..d4bd00a 100644 --- a/src/scene.cpp +++ b/src/scene.cpp @@ -48,10 +48,19 @@ int Scene::loadGeom(string objectid) { if (strcmp(line.c_str(), "sphere") == 0) { cout << "Creating new sphere..." << endl; newGeom.type = SPHERE; - } else if (strcmp(line.c_str(), "cube") == 0) { + } + else if (strcmp(line.c_str(), "cube") == 0) { cout << "Creating new cube..." << endl; newGeom.type = CUBE; } + else if (strcmp(line.c_str(), "csg1") == 0) { + cout << "Creating new csg1..." << endl; + newGeom.type = CSG1; + } + else if (strcmp(line.c_str(), "csg2") == 0) { + cout << "Creating new csg2..." << endl; + newGeom.type = CSG2; + } } //link material diff --git a/src/sceneStructs.h b/src/sceneStructs.h index b38b820..683ff1b 100644 --- a/src/sceneStructs.h +++ b/src/sceneStructs.h @@ -10,6 +10,8 @@ enum GeomType { SPHERE, CUBE, + CSG1, + CSG2 }; struct Ray {