From 30ea0c865d79504a836209e9be7bf5de79e96258 Mon Sep 17 00:00:00 2001 From: Ingo Wald Date: Fri, 16 Nov 2018 19:51:35 -0700 Subject: [PATCH] initial import --- .gitignore | 2 + CMakeLists.txt | 47 +++++ FinalChapter/CMakeLists.txt | 56 ++++++ FinalChapter/Camera.h | 31 +++ FinalChapter/Material.h | 103 ++++++++++ FinalChapter/finalChapter.cpp | 300 ++++++++++++++++++++++++++++ FinalChapter/programs/CuRandState.h | 38 ++++ FinalChapter/programs/dielectric.cu | 98 +++++++++ FinalChapter/programs/lambertian.cu | 69 +++++++ FinalChapter/programs/material.h | 56 ++++++ FinalChapter/programs/metal.cu | 70 +++++++ FinalChapter/programs/miss.cu | 35 ++++ FinalChapter/programs/prd.h | 32 +++ FinalChapter/programs/raygen.cu | 95 +++++++++ FinalChapter/programs/sampling.h | 37 ++++ FinalChapter/programs/sphere.cu | 78 ++++++++ FinalChapter/programs/vec.h | 117 +++++++++++ FinalChapter/savePPM.h | 45 +++++ LICENSE.txt | 201 +++++++++++++++++++ README.md | 1 + cmake/FindOptiX.cmake | 176 ++++++++++++++++ cmake/configure_optix.cmake | 68 +++++++ 22 files changed, 1755 insertions(+) create mode 100644 .gitignore create mode 100644 CMakeLists.txt create mode 100644 FinalChapter/CMakeLists.txt create mode 100644 FinalChapter/Camera.h create mode 100644 FinalChapter/Material.h create mode 100644 FinalChapter/finalChapter.cpp create mode 100644 FinalChapter/programs/CuRandState.h create mode 100644 FinalChapter/programs/dielectric.cu create mode 100644 FinalChapter/programs/lambertian.cu create mode 100644 FinalChapter/programs/material.h create mode 100644 FinalChapter/programs/metal.cu create mode 100644 FinalChapter/programs/miss.cu create mode 100644 FinalChapter/programs/prd.h create mode 100644 FinalChapter/programs/raygen.cu create mode 100644 FinalChapter/programs/sampling.h create mode 100644 FinalChapter/programs/sphere.cu create mode 100644 FinalChapter/programs/vec.h create mode 100644 FinalChapter/savePPM.h create mode 100644 LICENSE.txt create mode 100644 README.md create mode 100644 cmake/FindOptiX.cmake create mode 100644 cmake/configure_optix.cmake diff --git a/.gitignore b/.gitignore new file mode 100644 index 0000000..1865b49 --- /dev/null +++ b/.gitignore @@ -0,0 +1,2 @@ +*~ + diff --git a/CMakeLists.txt b/CMakeLists.txt new file mode 100644 index 0000000..55177d3 --- /dev/null +++ b/CMakeLists.txt @@ -0,0 +1,47 @@ +# ======================================================================== # +# Copyright 2018 Ingo Wald # +# # +# Licensed under the Apache License, Version 2.0 (the "License"); # +# you may not use this file except in compliance with the License. # +# You may obtain a copy of the License at # +# # +# http://www.apache.org/licenses/LICENSE-2.0 # +# # +# Unless required by applicable law or agreed to in writing, software # +# distributed under the License is distributed on an "AS IS" BASIS, # +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. # +# See the License for the specific language governing permissions and # +# limitations under the License. # +# ======================================================================== # + +project(RTOW-OptiX) + +cmake_minimum_required(VERSION 2.8) + +include(cmake/configure_optix.cmake) + + +if(NOT SET_UP_CONFIGURATIONS_DONE) + set(SET_UP_CONFIGURATIONS_DONE 1) + + # No reason to set CMAKE_CONFIGURATION_TYPES if it's not a multiconfig generator + # Also no reason mess with CMAKE_BUILD_TYPE if it's a multiconfig generator. + if(CMAKE_CONFIGURATION_TYPES) # multiconfig generator? + set(CMAKE_CONFIGURATION_TYPES "Debug;Release;Profile" CACHE STRING "" FORCE) + else() + if(NOT CMAKE_BUILD_TYPE) +# message("Defaulting to release build.") + set(CMAKE_BUILD_TYPE Release CACHE STRING "" FORCE) + endif() + set_property(CACHE CMAKE_BUILD_TYPE PROPERTY HELPSTRING "Choose the type of build") + # set the valid options for cmake-gui drop-down list + set_property(CACHE CMAKE_BUILD_TYPE PROPERTY STRINGS "Debug;Release;Profile") + endif() +endif() +mark_as_advanced(CUDA_SDK_ROOT_DIR) + +SET(CMAKE_RUNTIME_OUTPUT_DIRECTORY ${PROJECT_BINARY_DIR}) +SET(CMAKE_LIBRARY_OUTPUT_DIRECTORY ${PROJECT_BINARY_DIR}) +SET(CMAKE_ARCHIVE_OUTPUT_DIRECTORY ${PROJECT_BINARY_DIR}) + +add_subdirectory(FinalChapter) diff --git a/FinalChapter/CMakeLists.txt b/FinalChapter/CMakeLists.txt new file mode 100644 index 0000000..042ac65 --- /dev/null +++ b/FinalChapter/CMakeLists.txt @@ -0,0 +1,56 @@ +# ======================================================================== # +# Copyright 2018 Ingo Wald # +# # +# Licensed under the Apache License, Version 2.0 (the "License"); # +# you may not use this file except in compliance with the License. # +# You may obtain a copy of the License at # +# # +# http://www.apache.org/licenses/LICENSE-2.0 # +# # +# Unless required by applicable law or agreed to in writing, software # +# distributed under the License is distributed on an "AS IS" BASIS, # +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. # +# See the License for the specific language governing permissions and # +# limitations under the License. # +# ======================================================================== # + +# do some cmake magic to pre-compile the cuda file to ptx, and embed +# the resulting ptx code string into the final executable +cuda_compile_and_embed( + embedded_raygen_program programs/raygen.cu + ) +cuda_compile_and_embed( + embedded_sphere_programs programs/sphere.cu + ) +cuda_compile_and_embed( + embedded_miss_program programs/miss.cu + ) +cuda_compile_and_embed( + embedded_metal_programs programs/metal.cu + ) +cuda_compile_and_embed( + embedded_dielectric_programs programs/dielectric.cu + ) +cuda_compile_and_embed( + embedded_lambertian_programs programs/lambertian.cu + ) + +# this is doing the same using OptiX +add_executable(finalChapter + # C++ host code + programs/vec.h + programs/material.h + finalChapter.cpp + # embedded cuda kernels: + ${embedded_raygen_program} + ${embedded_sphere_programs} + ${embedded_lambertian_programs} + ${embedded_miss_program} + ${embedded_metal_programs} + ${embedded_dielectric_programs} + ) + +target_link_libraries(finalChapter + ${optix_LIBRARY} + ) + diff --git a/FinalChapter/Camera.h b/FinalChapter/Camera.h new file mode 100644 index 0000000..810f7f2 --- /dev/null +++ b/FinalChapter/Camera.h @@ -0,0 +1,31 @@ +__device__ vec3 random_in_unit_disk(curandState *local_rand_state) { + vec3 p; + do { + p = 2.0f*vec3(curand_uniform(local_rand_state), curand_uniform(local_rand_state), 0) - vec3(1, 1, 0); + } while (dot(p, p) >= 1.0f); + return p; +} + +class camera { +public: + __device__ camera(vec3 lookfrom, vec3 lookat, vec3 vup, float vfov, float aspect, float aperture, float focus_dist) { // vfov is top to bottom in degrees + lens_radius = aperture / 2.0f; + float theta = vfov * ((float)M_PI) / 180.0f; + float half_height = tan(theta / 2.0f); + float half_width = aspect * half_height; + origin = lookfrom; + w = unit_vector(lookfrom - lookat); + u = unit_vector(cross(vup, w)); + v = cross(w, u); + lower_left_corner = origin - half_width * focus_dist*u - half_height * focus_dist*v - focus_dist * w; + horizontal = 2.0f*half_width*focus_dist*u; + vertical = 2.0f*half_height*focus_dist*v; + } + + vec3 origin; + vec3 lower_left_corner; + vec3 horizontal; + vec3 vertical; + vec3 u, v, w; + float lens_radius; +}; diff --git a/FinalChapter/Material.h b/FinalChapter/Material.h new file mode 100644 index 0000000..f7473db --- /dev/null +++ b/FinalChapter/Material.h @@ -0,0 +1,103 @@ + +__device__ float schlick(float cosine, float ref_idx) { + float r0 = (1.0f - ref_idx) / (1.0f + ref_idx); + r0 = r0 * r0; + return r0 + (1.0f - r0)*pow((1.0f - cosine), 5.0f); +} + +__device__ bool refract(const vec3& v, const vec3& n, float ni_over_nt, vec3& refracted) { + vec3 uv = unit_vector(v); + float dt = dot(uv, n); + float discriminant = 1.0f - ni_over_nt * ni_over_nt*(1 - dt * dt); + if (discriminant > 0) { + refracted = ni_over_nt * (uv - n * dt) - n * sqrt(discriminant); + return true; + } + else + return false; +} + +#define RANDVEC3 vec3(curand_uniform(local_rand_state),curand_uniform(local_rand_state),curand_uniform(local_rand_state)) + +__device__ vec3 random_in_unit_sphere(curandState *local_rand_state) { + vec3 p; + do { + p = 2.0f*RANDVEC3 - vec3(1, 1, 1); + } while (p.squared_length() >= 1.0f); + return p; +} + +__device__ vec3 reflect(const vec3& v, const vec3& n) { + return v - 2.0f*dot(v, n)*n; +} + +class material { +public: + __device__ virtual bool scatter(const ray& r_in, const hit_record& rec, vec3& attenuation, ray& scattered, curandState *local_rand_state) const = 0; +}; + +class lambertian : public material { +public: + __device__ lambertian(const vec3& a) : albedo(a) {} + __device__ virtual bool scatter(const ray& r_in, const hit_record& rec, vec3& attenuation, ray& scattered, curandState *local_rand_state) const { + vec3 target = rec.p + rec.normal + random_in_unit_sphere(local_rand_state); + scattered = ray(rec.p, target - rec.p); + attenuation = albedo; + return true; + } + + vec3 albedo; +}; + +class metal : public material { +public: + __device__ metal(const vec3& a, float f) : albedo(a) { if (f < 1) fuzz = f; else fuzz = 1; } + __device__ virtual bool scatter(const ray& r_in, const hit_record& rec, vec3& attenuation, ray& scattered, curandState *local_rand_state) const { + vec3 reflected = reflect(unit_vector(r_in.direction()), rec.normal); + scattered = ray(rec.p, reflected + fuzz * random_in_unit_sphere(local_rand_state)); + attenuation = albedo; + return (dot(scattered.direction(), rec.normal) > 0.0f); + } + vec3 albedo; + float fuzz; +}; + +class dielectric : public material { +public: + __device__ dielectric(float ri) : ref_idx(ri) {} + __device__ virtual bool scatter(const ray& r_in, + const hit_record& rec, + vec3& attenuation, + ray& scattered, + curandState *local_rand_state) const { + vec3 outward_normal; + vec3 reflected = reflect(r_in.direction(), rec.normal); + float ni_over_nt; + attenuation = vec3(1.0, 1.0, 1.0); + vec3 refracted; + float reflect_prob; + float cosine; + if (dot(r_in.direction(), rec.normal) > 0.0f) { + outward_normal = -rec.normal; + ni_over_nt = ref_idx; + cosine = dot(r_in.direction(), rec.normal) / r_in.direction().length(); + cosine = sqrt(1.0f - ref_idx * ref_idx*(1 - cosine * cosine)); + } + else { + outward_normal = rec.normal; + ni_over_nt = 1.0f / ref_idx; + cosine = -dot(r_in.direction(), rec.normal) / r_in.direction().length(); + } + if (refract(r_in.direction(), outward_normal, ni_over_nt, refracted)) + reflect_prob = schlick(cosine, ref_idx); + else + reflect_prob = 1.0f; + if (curand_uniform(local_rand_state) < reflect_prob) + scattered = ray(rec.p, reflected); + else + scattered = ray(rec.p, refracted); + return true; + } + + float ref_idx; +}; diff --git a/FinalChapter/finalChapter.cpp b/FinalChapter/finalChapter.cpp new file mode 100644 index 0000000..5263067 --- /dev/null +++ b/FinalChapter/finalChapter.cpp @@ -0,0 +1,300 @@ +// ======================================================================== // +// Copyright 2018 Ingo Wald // +// // +// Licensed under the Apache License, Version 2.0 (the "License"); // +// you may not use this file except in compliance with the License. // +// You may obtain a copy of the License at // +// // +// http://www.apache.org/licenses/LICENSE-2.0 // +// // +// Unless required by applicable law or agreed to in writing, software // +// distributed under the License is distributed on an "AS IS" BASIS, // +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. // +// See the License for the specific language governing permissions and // +// limitations under the License. // +// ======================================================================== // + +// ooawe +#include "programs/vec.h" +#include "savePPM.h" +// optix +#include +#include +// std +#define _USE_MATH_DEFINES 1 +#include +#include +#include +#include +#include +#include + +optix::Context g_context; + +/*! the precompiled programs/raygen.cu code (in ptx) that our + cmake magic will precompile (to ptx) and link to the generated + executable (ie, we can simply declare and usethis here as + 'extern'. */ +extern "C" const char embedded_sphere_programs[]; +extern "C" const char embedded_raygen_program[]; +extern "C" const char embedded_miss_program[]; +extern "C" const char embedded_metal_programs[]; +extern "C" const char embedded_dielectric_programs[]; +extern "C" const char embedded_lambertian_programs[]; + +float rnd() +{ + // static std::random_device rd; //Will be used to obtain a seed for the random number engine + static std::mt19937 gen(0); //Standard mersenne_twister_engine seeded with rd() + static std::uniform_real_distribution dis(0.f, 1.f); + return dis(gen); +} + +/*! abstraction for a material that can create, and parameterize, + a newly created GI's material and closest hit program */ +struct Material { + virtual void assignTo(optix::GeometryInstance gi) const = 0; +}; + +/*! host side code for the "Lambertian" material; the actual + sampling code is in the programs/lambertian.cu closest hit program */ +struct Lambertian : public Material { + /*! constructor */ + Lambertian(const vec3f &albedo) : albedo(albedo) {} + /* create optix material, and assign mat and mat values to geom instance */ + virtual void assignTo(optix::GeometryInstance gi) const override { + optix::Material mat = g_context->createMaterial(); + mat->setClosestHitProgram(0, g_context->createProgramFromPTXString + (embedded_lambertian_programs, + "closest_hit")); + gi->setMaterial(/*ray type:*/0, mat); + gi["albedo"]->set3fv(&albedo.x); + } + const vec3f albedo; +}; + +/*! host side code for the "Metal" material; the actual + sampling code is in the programs/metal.cu closest hit program */ +struct Metal : public Material { + /*! constructor */ + Metal(const vec3f &albedo, const float fuzz) : albedo(albedo), fuzz(fuzz) {} + /* create optix material, and assign mat and mat values to geom instance */ + virtual void assignTo(optix::GeometryInstance gi) const override { + optix::Material mat = g_context->createMaterial(); + mat->setClosestHitProgram(0, g_context->createProgramFromPTXString + (embedded_metal_programs, + "closest_hit")); + gi->setMaterial(/*ray type:*/0, mat); + gi["albedo"]->set3fv(&albedo.x); + gi["fuzz"]->setFloat(fuzz); + } + const vec3f albedo; + const float fuzz; +}; + +/*! host side code for the "Dielectric" material; the actual + sampling code is in the programs/dielectric.cu closest hit program */ +struct Dielectric : public Material { + /*! constructor */ + Dielectric(const float ref_idx) : ref_idx(ref_idx) {} + /* create optix material, and assign mat and mat values to geom instance */ + virtual void assignTo(optix::GeometryInstance gi) const override { + optix::Material mat = g_context->createMaterial(); + mat->setClosestHitProgram(0, g_context->createProgramFromPTXString + (embedded_dielectric_programs, + "closest_hit")); + gi->setMaterial(/*ray type:*/0, mat); + gi["ref_idx"]->setFloat(ref_idx); + } + const float ref_idx; +}; + +optix::GeometryInstance createSphere(const vec3f ¢er, const float radius, const Material &material) +{ + optix::Geometry geometry = g_context->createGeometry(); + geometry->setPrimitiveCount(1); + geometry->setBoundingBoxProgram + (g_context->createProgramFromPTXString(embedded_sphere_programs, "get_bounds")); + geometry->setIntersectionProgram + (g_context->createProgramFromPTXString(embedded_sphere_programs, "hit_sphere")); + geometry["center"]->setFloat(center.x,center.y,center.z); + geometry["radius"]->setFloat(radius); + + optix::GeometryInstance gi = g_context->createGeometryInstance(); + gi->setGeometry(geometry); + gi->setMaterialCount(1); + material.assignTo(gi); + return gi; +} + +optix::GeometryGroup createScene() +{ + // first, create all geometry instances (GIs), and, for now, + // store them in a std::vector. For ease of reference, I'll + // stick wit the 'd_list' and 'd_world' names used in the + // reference C++ and CUDA codes. + std::vector d_list; + + d_list.push_back(createSphere(vec3f(0.f, -1000.0f, -1.f), 1000.f, + Lambertian(vec3f(0.5f, 0.5f, 0.5f)))); + + for (int a = -11; a < 11; a++) { + for (int b = -11; b < 11; b++) { + float choose_mat = rnd(); + vec3f center(a + rnd(), 0.2f, b + rnd()); + if (choose_mat < 0.8f) { + d_list.push_back(createSphere(center, 0.2f, + Lambertian(vec3f(rnd()*rnd(), rnd()*rnd(), rnd()*rnd())))); + } + else if (choose_mat < 0.95f) { + d_list.push_back(createSphere(center, 0.2f, + Metal(vec3f(0.5f*(1.0f + rnd()), 0.5f*(1.0f + rnd()), 0.5f*(1.0f + rnd())), 0.5f*rnd()))); + } + else { + d_list.push_back(createSphere(center, 0.2f, Dielectric(1.5f))); + } + } + } + d_list.push_back(createSphere(vec3f(0.f, 1.f, 0.f), 1.f, Dielectric(1.5f))); + d_list.push_back(createSphere(vec3f(-4.f, 1.f, 0.f), 1.f, Lambertian(vec3f(0.4f, 0.2f, 0.1f)))); + d_list.push_back(createSphere(vec3f(4.f, 1.f, 0.f), 1.f, Metal(vec3f(0.7f, 0.6f, 0.5f), 0.0f))); + + // now, create the optix world that contains all these GIs + optix::GeometryGroup d_world = g_context->createGeometryGroup(); + d_world->setAcceleration(g_context->createAcceleration("Bvh")); + d_world->setChildCount((int)d_list.size()); + for (int i = 0; i < d_list.size(); i++) + d_world->setChild(i, d_list[i]); + + // that all we have to do, the rest is up to optix + return d_world; +} + +struct Camera { + Camera(const vec3f &lookfrom, const vec3f &lookat, const vec3f &vup, + float vfov, float aspect, float aperture, float focus_dist) + { // vfov is top to bottom in degrees + lens_radius = aperture / 2.0f; + float theta = vfov * ((float)M_PI) / 180.0f; + float half_height = tan(theta / 2.0f); + float half_width = aspect * half_height; + origin = lookfrom; + w = unit_vector(lookfrom - lookat); + u = unit_vector(cross(vup, w)); + v = cross(w, u); + lower_left_corner = origin - half_width * focus_dist*u - half_height * focus_dist*v - focus_dist * w; + horizontal = 2.0f*half_width*focus_dist*u; + vertical = 2.0f*half_height*focus_dist*v; + } + + void set() + { + g_context["camera_lower_left_corner"]->set3fv(&lower_left_corner.x); + g_context["camera_horizontal"]->set3fv(&horizontal.x); + g_context["camera_vertical"]->set3fv(&vertical.x); + g_context["camera_origin"]->set3fv(&origin.x); + g_context["camera_u"]->set3fv(&u.x); + g_context["camera_v"]->set3fv(&v.x); + g_context["camera_w"]->set3fv(&w.x); + g_context["camera_lens_radius"]->setFloat(lens_radius); + } + vec3f origin; + vec3f lower_left_corner; + vec3f horizontal; + vec3f vertical; + vec3f u, v, w; + float lens_radius; +}; + + +void renderFrame(int Nx, int Ny) +{ + // ... and validate everything before launch. + g_context->validate(); + + // now that everything is set up: launch that ray generation program + g_context->launch(/*program ID:*/0, + /*launch dimensions:*/Nx, Ny); +} + +optix::Buffer createFrameBuffer(int Nx, int Ny) +{ + // ... create an image - as a 2D buffer of float3's ... + optix::Buffer pixelBuffer + = g_context->createBuffer(RT_BUFFER_OUTPUT); + pixelBuffer->setFormat(RT_FORMAT_FLOAT3); + pixelBuffer->setSize(Nx, Ny); + return pixelBuffer; +} + +void setRayGenProgram() +{ + optix::Program rayGenAndBackgroundProgram + = g_context->createProgramFromPTXString(embedded_raygen_program, + "renderPixel"); + g_context->setEntryPointCount(1); + g_context->setRayGenerationProgram(/*program ID:*/0, rayGenAndBackgroundProgram); +} + +void setMissProgram() +{ + optix::Program missProgram + = g_context->createProgramFromPTXString(embedded_miss_program, + "miss_program"); + g_context->setMissProgram(/*program ID:*/0, missProgram); +} + +int main(int ac, char **av) +{ + // before doing anything else: create a optix context + g_context = optix::Context::create(); + g_context->setRayTypeCount(1); + g_context->setStackSize( 8000 ); + + // define some image size ... + const size_t Nx = 1200, Ny = 800; + + // create - and set - the camera + const vec3f lookfrom(13, 2, 3); + const vec3f lookat(0, 0, 0); + Camera camera(lookfrom, + lookat, + /* up */ vec3f(0, 1, 0), + /* fovy, in degrees */ 20.0, + /* aspect */ float(Nx) / float(Ny), + /* aperture */ 0.1f, + /* dist to focus: */ 10.f); + camera.set(); + + // set the ray generation and miss shader program + setRayGenProgram(); + setMissProgram(); + + // create a frame buffer + optix::Buffer fb = createFrameBuffer(Nx, Ny); + g_context["fb"]->set(fb); + + // create the world to render + optix::GeometryGroup world = createScene(); + g_context["world"]->set(world); + + const int numSamples = 128; + g_context["numSamples"]->setInt(numSamples); + + // render the frame (and time it) + auto t0 = std::chrono::system_clock::now(); + renderFrame(Nx, Ny); + auto t1 = std::chrono::system_clock::now(); + std::cout << "done rendering, which took " + << std::setprecision(2) << std::chrono::duration(t1-t0).count() + << " seconds (for " << numSamples << " paths per pixel)" << std::endl; + + // ... map it, save it, and cleanly unmap it after reading... + const vec3f *pixels = (const vec3f *)fb->map(); + savePPM("finalChapter.ppm",Nx,Ny,pixels); + fb->unmap(); + + // ... done. + return 0; +} + diff --git a/FinalChapter/programs/CuRandState.h b/FinalChapter/programs/CuRandState.h new file mode 100644 index 0000000..2843e58 --- /dev/null +++ b/FinalChapter/programs/CuRandState.h @@ -0,0 +1,38 @@ +// ======================================================================== // +// Copyright 2018 Ingo Wald // +// // +// Licensed under the Apache License, Version 2.0 (the "License"); // +// you may not use this file except in compliance with the License. // +// You may obtain a copy of the License at // +// // +// http://www.apache.org/licenses/LICENSE-2.0 // +// // +// Unless required by applicable law or agreed to in writing, software // +// distributed under the License is distributed on an "AS IS" BASIS, // +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. // +// See the License for the specific language governing permissions and // +// limitations under the License. // +// ======================================================================== // + +#pragma once + +#include "vec.h" +#include + +struct CuRandState +{ + /*! initialize the random number generator with a new seed (usually + per pixel) */ + inline __device__ void init(int seed = 0) + { + curand_init(1984, seed, 0, &state); + } + + /*! get the next 'random' number in the sequence */ + inline __device__ float operator() () + { + return curand_uniform(&state); + } + + curandState state; +}; diff --git a/FinalChapter/programs/dielectric.cu b/FinalChapter/programs/dielectric.cu new file mode 100644 index 0000000..ac5ebbc --- /dev/null +++ b/FinalChapter/programs/dielectric.cu @@ -0,0 +1,98 @@ +// ======================================================================== // +// Copyright 2018 Ingo Wald // +// // +// Licensed under the Apache License, Version 2.0 (the "License"); // +// you may not use this file except in compliance with the License. // +// You may obtain a copy of the License at // +// // +// http://www.apache.org/licenses/LICENSE-2.0 // +// // +// Unless required by applicable law or agreed to in writing, software // +// distributed under the License is distributed on an "AS IS" BASIS, // +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. // +// See the License for the specific language governing permissions and // +// limitations under the License. // +// ======================================================================== // + +#include "material.h" +#include "prd.h" +#include "sampling.h" + +/*! the implicit state's ray we will intersect against */ +rtDeclareVariable(optix::Ray, ray, rtCurrentRay, ); +/*! the per ray data we operate on */ +rtDeclareVariable(PerRayData, prd, rtPayload, ); +rtDeclareVariable(rtObject, world, , ); + + +/*! the attributes we use to communicate between intersection programs and hit program */ +rtDeclareVariable(float3, hit_rec_normal, attribute hit_rec_normal, ); +rtDeclareVariable(float3, hit_rec_p, attribute hit_rec_p, ); + + +/*! and finally - that particular material's parameters */ +rtDeclareVariable(float, ref_idx, , ); + + + +/*! the actual scatter function - in Pete's reference code, that's a + virtual function, but since we have a different function per program + we do not need this here */ +inline __device__ bool scatter(const optix::Ray &r_in, + vec3f &attenuation, + optix::Ray &scattered + ) +{ + vec3f outward_normal; + vec3f reflected = reflect(r_in.direction, hit_rec_normal); + float ni_over_nt; + attenuation = vec3f(1.f, 1.f, 1.f); + vec3f refracted; + float reflect_prob; + float cosine; + CuRandState &rnd = *prd.cuRandState; + + if (dot(r_in.direction, hit_rec_normal) > 0.f) { + outward_normal = -hit_rec_normal; + ni_over_nt = ref_idx; + cosine = dot(r_in.direction, hit_rec_normal) / vec3f(r_in.direction).length(); + cosine = sqrtf(1.f - ref_idx*ref_idx*(1.f-cosine*cosine)); + } + else { + outward_normal = hit_rec_normal; + ni_over_nt = 1.0 / ref_idx; + cosine = -dot(r_in.direction, hit_rec_normal) / vec3f(r_in.direction).length(); + } + if (refract(r_in.direction, outward_normal, ni_over_nt, refracted)) + reflect_prob = schlick(cosine, ref_idx); + else + reflect_prob = 1.f; + if (rnd() < reflect_prob) + scattered = optix::Ray(/*org */hit_rec_p, + /*dir */reflected.as_float3(), + /*type*/0, + /*tmin*/1e-3f, + /*tmax*/RT_DEFAULT_MAX); + else + scattered = optix::Ray(/*org */hit_rec_p, + /*dir */refracted.as_float3(), + /*type*/0, + /*tmin*/1e-3f, + /*tmax*/RT_DEFAULT_MAX); + return true; +} + +RT_PROGRAM void closest_hit() +{ + optix::Ray scattered; + vec3f attenuation; + if (prd.depth < 50 && scatter(ray,attenuation,scattered)) { + PerRayData rec; + rec.depth = prd.depth+1; + rec.cuRandState = prd.cuRandState; + rtTrace(world,scattered,rec); + prd.color = attenuation * rec.color; + } else { + prd.color = vec3f(0,0,0); + } +} diff --git a/FinalChapter/programs/lambertian.cu b/FinalChapter/programs/lambertian.cu new file mode 100644 index 0000000..8bdbf65 --- /dev/null +++ b/FinalChapter/programs/lambertian.cu @@ -0,0 +1,69 @@ +// ======================================================================== // +// Copyright 2018 Ingo Wald // +// // +// Licensed under the Apache License, Version 2.0 (the "License"); // +// you may not use this file except in compliance with the License. // +// You may obtain a copy of the License at // +// // +// http://www.apache.org/licenses/LICENSE-2.0 // +// // +// Unless required by applicable law or agreed to in writing, software // +// distributed under the License is distributed on an "AS IS" BASIS, // +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. // +// See the License for the specific language governing permissions and // +// limitations under the License. // +// ======================================================================== // + +#include "material.h" +#include "prd.h" +#include "sampling.h" + +/*! the implicit state's ray we will intersect against */ +rtDeclareVariable(optix::Ray, ray, rtCurrentRay, ); +/*! the per ray data we operate on */ +rtDeclareVariable(PerRayData, prd, rtPayload, ); +rtDeclareVariable(rtObject, world, , ); + + + +/*! the attributes we use to communicate between intersection programs and hit program */ +rtDeclareVariable(float3, hit_rec_normal, attribute hit_rec_normal, ); +rtDeclareVariable(float3, hit_rec_p, attribute hit_rec_p, ); + + +/*! and finally - that particular material's parameters */ +rtDeclareVariable(float3, albedo, , ); + + + +/*! the actual scatter function - in Pete's reference code, that's a + virtual function, but since we have a different function per program + we do not need this here */ +inline __device__ bool scatter(const optix::Ray &ray_in, + vec3f &attenuation, + optix::Ray &scattered) +{ + vec3f target = hit_rec_p + hit_rec_normal + random_in_unit_sphere(*prd.cuRandState); + scattered = optix::Ray(/*org */hit_rec_p, + /*dir */(target-hit_rec_p).as_float3(), + /*type*/0, + /*tmin*/1e-3f, + /*tmax*/RT_DEFAULT_MAX); + attenuation = albedo; + return true; +} + +RT_PROGRAM void closest_hit() +{ + optix::Ray scattered; + vec3f attenuation; + if (prd.depth < 50 && scatter(ray,attenuation,scattered)) { + PerRayData rec; + rec.depth = prd.depth+1; + rec.cuRandState = prd.cuRandState; + rtTrace(world,scattered,rec); + prd.color = attenuation * rec.color; + } else { + prd.color = vec3f(0,0,0); + } +} diff --git a/FinalChapter/programs/material.h b/FinalChapter/programs/material.h new file mode 100644 index 0000000..e4b1320 --- /dev/null +++ b/FinalChapter/programs/material.h @@ -0,0 +1,56 @@ +// ======================================================================== // +// Copyright 2018 Ingo Wald // +// // +// Licensed under the Apache License, Version 2.0 (the "License"); // +// you may not use this file except in compliance with the License. // +// You may obtain a copy of the License at // +// // +// http://www.apache.org/licenses/LICENSE-2.0 // +// // +// Unless required by applicable law or agreed to in writing, software // +// distributed under the License is distributed on an "AS IS" BASIS, // +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. // +// See the License for the specific language governing permissions and // +// limitations under the License. // +// ======================================================================== // + +#pragma once + +#include "prd.h" +#include "CuRandState.h" + + +__device__ float schlick(float cosine, float ref_idx) { + float r0 = (1.0f - ref_idx) / (1.0f + ref_idx); + r0 = r0 * r0; + return r0 + (1.0f - r0)*pow((1.0f - cosine), 5.0f); +} + +__device__ bool refract(const vec3f& v, const vec3f& n, float ni_over_nt, vec3f& refracted) { + vec3f uv = unit_vector(v); + float dt = dot(uv, n); + float discriminant = 1.0f - ni_over_nt * ni_over_nt*(1 - dt * dt); + if (discriminant > 0) { + refracted = ni_over_nt * (uv - n * dt) - n * sqrt(discriminant); + return true; + } + else + return false; +} + +#define RANDVEC3F vec3f(curand_uniform(local_rand_state),curand_uniform(local_rand_state),curand_uniform(local_rand_state)) + +__device__ vec3f random_in_unit_sphere(curandState *local_rand_state) { + vec3f p; + do { + p = 2.0f*RANDVEC3F - vec3f(1, 1, 1); + } while (p.squared_length() >= 1.0f); + return p; +} + + +inline __device__ vec3f reflect(const vec3f &v, const vec3f &n) +{ + return v - 2.0f*dot(v, n)*n; +} + diff --git a/FinalChapter/programs/metal.cu b/FinalChapter/programs/metal.cu new file mode 100644 index 0000000..e74eac9 --- /dev/null +++ b/FinalChapter/programs/metal.cu @@ -0,0 +1,70 @@ +// ======================================================================== // +// Copyright 2018 Ingo Wald // +// // +// Licensed under the Apache License, Version 2.0 (the "License"); // +// you may not use this file except in compliance with the License. // +// You may obtain a copy of the License at // +// // +// http://www.apache.org/licenses/LICENSE-2.0 // +// // +// Unless required by applicable law or agreed to in writing, software // +// distributed under the License is distributed on an "AS IS" BASIS, // +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. // +// See the License for the specific language governing permissions and // +// limitations under the License. // +// ======================================================================== // + +#include "material.h" +#include "prd.h" +#include "sampling.h" + +/*! the implicit state's ray we will intersect against */ +rtDeclareVariable(optix::Ray, ray, rtCurrentRay, ); +/*! the per ray data we operate on */ +rtDeclareVariable(PerRayData, prd, rtPayload, ); +rtDeclareVariable(rtObject, world, , ); + + + +/*! the attributes we use to communicate between intersection programs and hit program */ +rtDeclareVariable(float3, hit_rec_normal, attribute hit_rec_normal, ); +rtDeclareVariable(float3, hit_rec_p, attribute hit_rec_p, ); + + +/*! and finally - that particular material's parameters */ +rtDeclareVariable(float3, albedo, , ); +rtDeclareVariable(float, fuzz, , ); + + +/*! the actual scatter function - in Pete's reference code, that's a + virtual function, but since we have a different function per program + we do not need this here */ +inline __device__ bool scatter(const optix::Ray &ray_in, + vec3f &attenuation, + optix::Ray &scattered) +{ + vec3f reflected = reflect(unit_vector(ray_in.direction),hit_rec_normal); + scattered = optix::Ray(/*org */hit_rec_p, + /*dir */(reflected+fuzz*random_in_unit_sphere(*prd.cuRandState)) + .as_float3(), + /*type*/0, + /*tmin*/1e-3f, + /*tmax*/RT_DEFAULT_MAX); + attenuation = albedo; + return (dot(scattered.direction, hit_rec_normal) > 0.f); +} + +RT_PROGRAM void closest_hit() +{ + optix::Ray scattered; + vec3f attenuation; + if (prd.depth < 50 && scatter(ray,attenuation,scattered)) { + PerRayData rec; + rec.depth = prd.depth+1; + rec.cuRandState = prd.cuRandState; + rtTrace(world,scattered,rec); + prd.color = attenuation * rec.color; + } else { + prd.color = vec3f(0,0,0); + } +} diff --git a/FinalChapter/programs/miss.cu b/FinalChapter/programs/miss.cu new file mode 100644 index 0000000..5250a50 --- /dev/null +++ b/FinalChapter/programs/miss.cu @@ -0,0 +1,35 @@ +// ======================================================================== // +// Copyright 2018 Ingo Wald // +// // +// Licensed under the Apache License, Version 2.0 (the "License"); // +// you may not use this file except in compliance with the License. // +// You may obtain a copy of the License at // +// // +// http://www.apache.org/licenses/LICENSE-2.0 // +// // +// Unless required by applicable law or agreed to in writing, software // +// distributed under the License is distributed on an "AS IS" BASIS, // +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. // +// See the License for the specific language governing permissions and // +// limitations under the License. // +// ======================================================================== // + +#include "material.h" + +/*! the implicit state's ray we will intersect against */ +rtDeclareVariable(optix::Ray, ray, rtCurrentRay, ); +/*! the per ray data we operate on */ +rtDeclareVariable(PerRayData, prd, rtPayload, ); + +inline __device__ vec3f missColor(const optix::Ray &ray) +{ + const vec3f unit_direction = normalize(ray.direction); + const float t = 0.5f*(unit_direction.y + 1.0f); + const vec3f c = (1.0f - t)*vec3f(1.0f, 1.0f, 1.0f) + t * vec3f(0.5f, 0.7f, 1.0f); + return c; +} + +RT_PROGRAM void miss_program() +{ + prd.color = missColor(ray); +} diff --git a/FinalChapter/programs/prd.h b/FinalChapter/programs/prd.h new file mode 100644 index 0000000..fd7971b --- /dev/null +++ b/FinalChapter/programs/prd.h @@ -0,0 +1,32 @@ +// ======================================================================== // +// Copyright 2018 Ingo Wald // +// // +// Licensed under the Apache License, Version 2.0 (the "License"); // +// you may not use this file except in compliance with the License. // +// You may obtain a copy of the License at // +// // +// http://www.apache.org/licenses/LICENSE-2.0 // +// // +// Unless required by applicable law or agreed to in writing, software // +// distributed under the License is distributed on an "AS IS" BASIS, // +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. // +// See the License for the specific language governing permissions and // +// limitations under the License. // +// ======================================================================== // + +#pragma once + +#include "vec.h" +#include "CuRandState.h" + +/*! "per ray data" (PRD) for our sample's rays. In the simple example, there is only + one ray type, and it only ever returns one thing, which is a color (everything else + is handled through the recursion). In addition to that return type, rays have to + carry recursion state, which in this case are recursion depth and random number state */ +struct PerRayData { + CuRandState *cuRandState; + /*! recursion depth */ + int depth; + /*! return value of the recursive trace() call */ + vec3f color; +}; diff --git a/FinalChapter/programs/raygen.cu b/FinalChapter/programs/raygen.cu new file mode 100644 index 0000000..080a765 --- /dev/null +++ b/FinalChapter/programs/raygen.cu @@ -0,0 +1,95 @@ +// ======================================================================== // +// Copyright 2018 Ingo Wald // +// // +// Licensed under the Apache License, Version 2.0 (the "License"); // +// you may not use this file except in compliance with the License. // +// You may obtain a copy of the License at // +// // +// http://www.apache.org/licenses/LICENSE-2.0 // +// // +// Unless required by applicable law or agreed to in writing, software // +// distributed under the License is distributed on an "AS IS" BASIS, // +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. // +// See the License for the specific language governing permissions and // +// limitations under the License. // +// ======================================================================== // + +// optix code: +#include +#include +#include "prd.h" +#include "sampling.h" + +/*! the 'builtin' launch index we need to render a frame */ +rtDeclareVariable(uint2, pixelID, rtLaunchIndex, ); +rtDeclareVariable(uint2, launchDim, rtLaunchDim, ); + +/*! the ray related state */ +rtDeclareVariable(optix::Ray, ray, rtCurrentRay, ); +rtDeclareVariable(PerRayData, prd, rtPayload, ); + +/*! the 2D, float3-type color frame buffer we'll write into */ +rtBuffer fb; + +rtDeclareVariable(int, numSamples, , ); + +rtDeclareVariable(rtObject, world, , ); + +rtDeclareVariable(float3, camera_lower_left_corner, , ); +rtDeclareVariable(float3, camera_horizontal, , ); +rtDeclareVariable(float3, camera_vertical, , ); +rtDeclareVariable(float3, camera_origin, , ); +rtDeclareVariable(float3, camera_u, , ); +rtDeclareVariable(float3, camera_v, , ); +rtDeclareVariable(float, camera_lens_radius, , ); + +struct Camera { + static __device__ optix::Ray generateRay(float s, float t, CuRandState &rnd) + { + const vec3f rd = camera_lens_radius * random_in_unit_disk(rnd); + const vec3f lens_offset = camera_u * rd.x + camera_v * rd.y; + const vec3f origin = camera_origin + lens_offset; + const vec3f direction + = camera_lower_left_corner + + s * camera_horizontal + + t * camera_vertical + - origin; + return optix::make_Ray(/* origin : */ origin.as_float3(), + /* direction: */ direction.as_float3(), + /* ray type : */ 0, + /* tmin : */ 1e-6f, + /* tmax : */ RT_DEFAULT_MAX); + } +}; + +inline __device__ vec3f color(optix::Ray ray, CuRandState &rnd) +{ + PerRayData prd; + prd.cuRandState = &rnd; + prd.depth = 0; + rtTrace(world, ray, prd); + return prd.color; +} + +inline __device__ float rnd() { return 0.5f; } + +/*! the actual ray generation program - note this has no formal + function parameters, but gets its paramters throught the 'pixelID' + and 'pixelBuffer' variables/buffers declared above */ +RT_PROGRAM void renderPixel() +{ + int pixel_index = pixelID.y * launchDim.x + pixelID.x; + vec3f col(0.f, 0.f, 0.f); + CuRandState rnd; + rnd.init(pixel_index); + for (int s = 0; s < numSamples; s++) { + float u = float(pixelID.x + rnd()) / float(launchDim.x); + float v = float(pixelID.y + rnd()) / float(launchDim.y); + optix::Ray ray = Camera::generateRay(u, v, rnd); + col += color(ray, rnd); + } + col = col / float(numSamples); + + fb[pixelID] = col.as_float3(); +} + diff --git a/FinalChapter/programs/sampling.h b/FinalChapter/programs/sampling.h new file mode 100644 index 0000000..37524b4 --- /dev/null +++ b/FinalChapter/programs/sampling.h @@ -0,0 +1,37 @@ +// ======================================================================== // +// Copyright 2018 Ingo Wald // +// // +// Licensed under the Apache License, Version 2.0 (the "License"); // +// you may not use this file except in compliance with the License. // +// You may obtain a copy of the License at // +// // +// http://www.apache.org/licenses/LICENSE-2.0 // +// // +// Unless required by applicable law or agreed to in writing, software // +// distributed under the License is distributed on an "AS IS" BASIS, // +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. // +// See the License for the specific language governing permissions and // +// limitations under the License. // +// ======================================================================== // + +#pragma once + +#include "vec.h" +#include "CuRandState.h" + +inline __device__ vec3f random_in_unit_disk(CuRandState &local_rand_state) { + vec3f p; + do { + p = 2.0f*vec3f(local_rand_state(), local_rand_state(), 0) - vec3f(1, 1, 0); + } while (dot(p, p) >= 1.0f); + return p; +} + +__device__ vec3f random_in_unit_sphere(CuRandState &local_rand_state) { + vec3f p; + do { + p = 2.0f*vec3f(local_rand_state(), local_rand_state(), local_rand_state()) - vec3f(1.f, 1.f, 1.f); + } while (p.squared_length() >= 1.0f); + return p; +} + diff --git a/FinalChapter/programs/sphere.cu b/FinalChapter/programs/sphere.cu new file mode 100644 index 0000000..1d3fdf1 --- /dev/null +++ b/FinalChapter/programs/sphere.cu @@ -0,0 +1,78 @@ +// ======================================================================== // +// Copyright 2018 Ingo Wald // +// // +// Licensed under the Apache License, Version 2.0 (the "License"); // +// you may not use this file except in compliance with the License. // +// You may obtain a copy of the License at // +// // +// http://www.apache.org/licenses/LICENSE-2.0 // +// // +// Unless required by applicable law or agreed to in writing, software // +// distributed under the License is distributed on an "AS IS" BASIS, // +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. // +// See the License for the specific language governing permissions and // +// limitations under the License. // +// ======================================================================== // + +#include +#include "prd.h" + +/*! the parameters that describe each individual sphere geometry */ +rtDeclareVariable(float3, center, , ); +rtDeclareVariable(float, radius, , ); + +/*! the implicit state's ray we will intersect against */ +rtDeclareVariable(optix::Ray, ray, rtCurrentRay, ); + +/*! the attributes we use to communicate between intersection programs and hit program */ +rtDeclareVariable(float3, hit_rec_normal, attribute hit_rec_normal, ); +rtDeclareVariable(float3, hit_rec_p, attribute hit_rec_p, ); + +/*! the per ray data we operate on */ +rtDeclareVariable(PerRayData, prd, rtPayload, ); + + +// Program that performs the ray-sphere intersection +// +// note that this is here is a simple, but not necessarily most numerically +// stable ray-sphere intersection variant out there. There are more +// stable variants out there, but for now let's stick with the one that +// the reference code used. +RT_PROGRAM void hit_sphere(int pid) +{ + const float3 oc = ray.origin - center; + const float a = dot(ray.direction, ray.direction); + const float b = dot(oc, ray.direction); + const float c = dot(oc, oc) - radius * radius; + const float discriminant = b * b - a * c; + + if (discriminant < 0.f) return; + + float temp = (-b - sqrtf(discriminant)) / a; + if (temp < ray.tmax && temp > ray.tmin) { + if (rtPotentialIntersection(temp)) { + hit_rec_p = ray.origin + temp * ray.direction; + hit_rec_normal = (hit_rec_p - center) / radius; + rtReportIntersection(0); + } + } + temp = (-b + sqrtf(discriminant)) / a; + if (temp < ray.tmax && temp > ray.tmin) { + if (rtPotentialIntersection(temp)) { + hit_rec_p = ray.origin + temp * ray.direction; + hit_rec_normal = (hit_rec_p - center) / radius; + rtReportIntersection(0); + } + } +} + +/*! returns the bounding box of the pid'th primitive + in this gometry. Since we only have one sphere in this + program (we handle multiple spheres by having a different + geometry per sphere), the'pid' parameter is ignored */ +RT_PROGRAM void get_bounds(int pid, float result[6]) +{ + optix::Aabb* aabb = (optix::Aabb*)result; + aabb->m_min = center - radius; + aabb->m_max = center + radius; +} diff --git a/FinalChapter/programs/vec.h b/FinalChapter/programs/vec.h new file mode 100644 index 0000000..0ec18d2 --- /dev/null +++ b/FinalChapter/programs/vec.h @@ -0,0 +1,117 @@ +// ======================================================================== // +// Copyright 2018 Ingo Wald // +// // +// Licensed under the Apache License, Version 2.0 (the "License"); // +// you may not use this file except in compliance with the License. // +// You may obtain a copy of the License at // +// // +// http://www.apache.org/licenses/LICENSE-2.0 // +// // +// Unless required by applicable law or agreed to in writing, software // +// distributed under the License is distributed on an "AS IS" BASIS, // +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. // +// See the License for the specific language governing permissions and // +// limitations under the License. // +// ======================================================================== // + +#pragma once + +// optix code: +#include +#include +#include + +#ifdef __CUDACC__ +#else +#endif + +struct vec3f +{ + inline __device__ vec3f() {} + inline __device__ vec3f(float f) + : x(f), y(f), z(f) + {} + inline __device__ vec3f(const float x, const float y, const float z) + : x(x), y(y), z(z) + {} + inline __device__ vec3f(const vec3f &v) + : x(v.x), y(v.y), z(v.z) + {} +#ifdef __CUDACC__ + inline __device__ vec3f(const float3 v) + : x(v.x), y(v.y), z(v.z) + {} + inline __device__ float3 as_float3() const { return make_float3(x, y, z); } +#endif + inline __device__ float squared_length() const { return x * x + y * y + z * z; } + inline __device__ float length() const { return sqrtf(x * x + y * y + z * z); } + float x, y, z; +}; + +inline __device__ vec3f operator*(const vec3f &a, const vec3f &b) +{ + return vec3f(a.x*b.x, a.y*b.y, a.z*b.z); +} + +inline __device__ vec3f operator/(const vec3f &a, const vec3f &b) +{ + return vec3f(a.x/b.x, a.y/b.y, a.z/b.z); +} + +inline __device__ vec3f operator-(const vec3f &a, const vec3f &b) +{ + return vec3f(a.x-b.x, a.y-b.y, a.z-b.z); +} + +inline __device__ vec3f operator-(const vec3f &a) +{ + return vec3f(-a.x, -a.y, -a.z); +} + +inline __device__ vec3f operator+(const vec3f &a, const vec3f &b) +{ + return vec3f(a.x+b.x, a.y+b.y, a.z+b.z); +} + +inline __device__ vec3f &operator+=(vec3f &a, const vec3f &b) +{ + a = a + b; + return a; +} + +inline __device__ float dot(const vec3f &a, const vec3f &b) +{ + return a.x*b.x + a.y*b.y + a.z*b.z; +} + +inline __device__ vec3f cross(const vec3f &a, const vec3f &b) +{ + return vec3f( + a.y*b.z - a.z*b.y, + a.z*b.x - a.x*b.z, + a.x*b.y - a.y*b.x); +} + +inline __device__ vec3f normalize(const vec3f &v) +{ + return v * (1.f / sqrtf(dot(v, v))); +} + + +inline __device__ vec3f unit_vector(const vec3f &v) +{ + return normalize(v); +} + +/*! return absolute value of each component */ +inline __device__ vec3f abs(const vec3f &v) +{ + return vec3f(fabsf(v.x),fabsf(v.y),fabsf(v.z)); +} + +inline __device__ vec3f mod(const vec3f &a,const vec3f &b) +{ + return vec3f(fmodf(a.x,b.x), + fmodf(a.y,b.y), + fmodf(a.z,b.z)); +} diff --git a/FinalChapter/savePPM.h b/FinalChapter/savePPM.h new file mode 100644 index 0000000..511acb8 --- /dev/null +++ b/FinalChapter/savePPM.h @@ -0,0 +1,45 @@ +// ======================================================================== // +// Copyright 2018 Ingo Wald // +// // +// Licensed under the Apache License, Version 2.0 (the "License"); // +// you may not use this file except in compliance with the License. // +// You may obtain a copy of the License at // +// // +// http://www.apache.org/licenses/LICENSE-2.0 // +// // +// Unless required by applicable law or agreed to in writing, software // +// distributed under the License is distributed on an "AS IS" BASIS, // +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. // +// See the License for the specific language governing permissions and // +// limitations under the License. // +// ======================================================================== // + +#pragma once + +// ooawe +#include "programs/vec.h" +// std +#include +#include +#include + +/*! saving to a 'P3' type PPM file (255 values per channel). There + are many other, more C++-like, ways of writing this; this version is + intentionally written as close to the RTOW version as possible */ +inline void savePPM(const std::string &fileName, + const size_t Nx, const size_t Ny, const vec3f *pixels) +{ + std::ofstream file(fileName); + assert(file.good()); + + file << "P3\n" << Nx << " " << Ny << "\n255\n"; + for (int iy=(int)Ny-1; iy>=0; --iy) + for (int ix=0; ix<(int)Nx; ix++) { + int ir = int(255.99*pixels[ix+Nx*iy].x); + int ig = int(255.99*pixels[ix+Nx*iy].y); + int ib = int(255.99*pixels[ix+Nx*iy].z); + file << ir << " " << ig << " " << ib << "\n"; + } + assert(file.good()); +} + diff --git a/LICENSE.txt b/LICENSE.txt new file mode 100644 index 0000000..6a25b75 --- /dev/null +++ b/LICENSE.txt @@ -0,0 +1,201 @@ + Apache License + Version 2.0, January 2004 + http://www.apache.org/licenses/ + + TERMS AND CONDITIONS FOR USE, REPRODUCTION, AND DISTRIBUTION + + 1. Definitions. + + "License" shall mean the terms and conditions for use, reproduction, + and distribution as defined by Sections 1 through 9 of this document. + + "Licensor" shall mean the copyright owner or entity authorized by + the copyright owner that is granting the License. + + "Legal Entity" shall mean the union of the acting entity and all + other entities that control, are controlled by, or are under common + control with that entity. For the purposes of this definition, + "control" means (i) the power, direct or indirect, to cause the + direction or management of such entity, whether by contract or + otherwise, or (ii) ownership of fifty percent (50%) or more of the + outstanding shares, or (iii) beneficial ownership of such entity. + + "You" (or "Your") shall mean an individual or Legal Entity + exercising permissions granted by this License. + + "Source" form shall mean the preferred form for making modifications, + including but not limited to software source code, documentation + source, and configuration files. + + "Object" form shall mean any form resulting from mechanical + transformation or translation of a Source form, including but + not limited to compiled object code, generated documentation, + and conversions to other media types. + + "Work" shall mean the work of authorship, whether in Source or + Object form, made available under the License, as indicated by a + copyright notice that is included in or attached to the work + (an example is provided in the Appendix below). + + "Derivative Works" shall mean any work, whether in Source or Object + form, that is based on (or derived from) the Work and for which the + editorial revisions, annotations, elaborations, or other modifications + represent, as a whole, an original work of authorship. For the purposes + of this License, Derivative Works shall not include works that remain + separable from, or merely link (or bind by name) to the interfaces of, + the Work and Derivative Works thereof. + + "Contribution" shall mean any work of authorship, including + the original version of the Work and any modifications or additions + to that Work or Derivative Works thereof, that is intentionally + submitted to Licensor for inclusion in the Work by the copyright owner + or by an individual or Legal Entity authorized to submit on behalf of + the copyright owner. For the purposes of this definition, "submitted" + means any form of electronic, verbal, or written communication sent + to the Licensor or its representatives, including but not limited to + communication on electronic mailing lists, source code control systems, + and issue tracking systems that are managed by, or on behalf of, the + Licensor for the purpose of discussing and improving the Work, but + excluding communication that is conspicuously marked or otherwise + designated in writing by the copyright owner as "Not a Contribution." + + "Contributor" shall mean Licensor and any individual or Legal Entity + on behalf of whom a Contribution has been received by Licensor and + subsequently incorporated within the Work. + + 2. Grant of Copyright License. Subject to the terms and conditions of + this License, each Contributor hereby grants to You a perpetual, + worldwide, non-exclusive, no-charge, royalty-free, irrevocable + copyright license to reproduce, prepare Derivative Works of, + publicly display, publicly perform, sublicense, and distribute the + Work and such Derivative Works in Source or Object form. + + 3. Grant of Patent License. Subject to the terms and conditions of + this License, each Contributor hereby grants to You a perpetual, + worldwide, non-exclusive, no-charge, royalty-free, irrevocable + (except as stated in this section) patent license to make, have made, + use, offer to sell, sell, import, and otherwise transfer the Work, + where such license applies only to those patent claims licensable + by such Contributor that are necessarily infringed by their + Contribution(s) alone or by combination of their Contribution(s) + with the Work to which such Contribution(s) was submitted. If You + institute patent litigation against any entity (including a + cross-claim or counterclaim in a lawsuit) alleging that the Work + or a Contribution incorporated within the Work constitutes direct + or contributory patent infringement, then any patent licenses + granted to You under this License for that Work shall terminate + as of the date such litigation is filed. + + 4. Redistribution. You may reproduce and distribute copies of the + Work or Derivative Works thereof in any medium, with or without + modifications, and in Source or Object form, provided that You + meet the following conditions: + + (a) You must give any other recipients of the Work or + Derivative Works a copy of this License; and + + (b) You must cause any modified files to carry prominent notices + stating that You changed the files; and + + (c) You must retain, in the Source form of any Derivative Works + that You distribute, all copyright, patent, trademark, and + attribution notices from the Source form of the Work, + excluding those notices that do not pertain to any part of + the Derivative Works; and + + (d) If the Work includes a "NOTICE" text file as part of its + distribution, then any Derivative Works that You distribute must + include a readable copy of the attribution notices contained + within such NOTICE file, excluding those notices that do not + pertain to any part of the Derivative Works, in at least one + of the following places: within a NOTICE text file distributed + as part of the Derivative Works; within the Source form or + documentation, if provided along with the Derivative Works; or, + within a display generated by the Derivative Works, if and + wherever such third-party notices normally appear. The contents + of the NOTICE file are for informational purposes only and + do not modify the License. You may add Your own attribution + notices within Derivative Works that You distribute, alongside + or as an addendum to the NOTICE text from the Work, provided + that such additional attribution notices cannot be construed + as modifying the License. + + You may add Your own copyright statement to Your modifications and + may provide additional or different license terms and conditions + for use, reproduction, or distribution of Your modifications, or + for any such Derivative Works as a whole, provided Your use, + reproduction, and distribution of the Work otherwise complies with + the conditions stated in this License. + + 5. Submission of Contributions. Unless You explicitly state otherwise, + any Contribution intentionally submitted for inclusion in the Work + by You to the Licensor shall be under the terms and conditions of + this License, without any additional terms or conditions. + Notwithstanding the above, nothing herein shall supersede or modify + the terms of any separate license agreement you may have executed + with Licensor regarding such Contributions. + + 6. Trademarks. This License does not grant permission to use the trade + names, trademarks, service marks, or product names of the Licensor, + except as required for reasonable and customary use in describing the + origin of the Work and reproducing the content of the NOTICE file. + + 7. Disclaimer of Warranty. Unless required by applicable law or + agreed to in writing, Licensor provides the Work (and each + Contributor provides its Contributions) on an "AS IS" BASIS, + WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or + implied, including, without limitation, any warranties or conditions + of TITLE, NON-INFRINGEMENT, MERCHANTABILITY, or FITNESS FOR A + PARTICULAR PURPOSE. You are solely responsible for determining the + appropriateness of using or redistributing the Work and assume any + risks associated with Your exercise of permissions under this License. + + 8. Limitation of Liability. In no event and under no legal theory, + whether in tort (including negligence), contract, or otherwise, + unless required by applicable law (such as deliberate and grossly + negligent acts) or agreed to in writing, shall any Contributor be + liable to You for damages, including any direct, indirect, special, + incidental, or consequential damages of any character arising as a + result of this License or out of the use or inability to use the + Work (including but not limited to damages for loss of goodwill, + work stoppage, computer failure or malfunction, or any and all + other commercial damages or losses), even if such Contributor + has been advised of the possibility of such damages. + + 9. Accepting Warranty or Additional Liability. While redistributing + the Work or Derivative Works thereof, You may choose to offer, + and charge a fee for, acceptance of support, warranty, indemnity, + or other liability obligations and/or rights consistent with this + License. However, in accepting such obligations, You may act only + on Your own behalf and on Your sole responsibility, not on behalf + of any other Contributor, and only if You agree to indemnify, + defend, and hold each Contributor harmless for any liability + incurred by, or claims asserted against, such Contributor by reason + of your accepting any such warranty or additional liability. + + END OF TERMS AND CONDITIONS + + APPENDIX: How to apply the Apache License to your work. + + To apply the Apache License to your work, attach the following + boilerplate notice, with the fields enclosed by brackets "{}" + replaced with your own identifying information. (Don't include + the brackets!) The text should be enclosed in the appropriate + comment syntax for the file format. We also recommend that a + file or class name and description of purpose be included on the + same "printed page" as the copyright notice for easier + identification within third-party archives. + + Copyright 2018 Ingo Wald + + Licensed under the Apache License, Version 2.0 (the "License"); + you may not use this file except in compliance with the License. + You may obtain a copy of the License at + + http://www.apache.org/licenses/LICENSE-2.0 + + Unless required by applicable law or agreed to in writing, software + distributed under the License is distributed on an "AS IS" BASIS, + WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + See the License for the specific language governing permissions and + limitations under the License. diff --git a/README.md b/README.md new file mode 100644 index 0000000..7cfdebe --- /dev/null +++ b/README.md @@ -0,0 +1 @@ +# RTOW-OptiX diff --git a/cmake/FindOptiX.cmake b/cmake/FindOptiX.cmake new file mode 100644 index 0000000..7c9430f --- /dev/null +++ b/cmake/FindOptiX.cmake @@ -0,0 +1,176 @@ +# +# Copyright (c) 2018 NVIDIA CORPORATION. All rights reserved. +# +# Redistribution and use in source and binary forms, with or without +# modification, are permitted provided that the following conditions +# are met: +# * Redistributions of source code must retain the above copyright +# notice, this list of conditions and the following disclaimer. +# * Redistributions in binary form must reproduce the above copyright +# notice, this list of conditions and the following disclaimer in the +# documentation and/or other materials provided with the distribution. +# * Neither the name of NVIDIA CORPORATION nor the names of its +# contributors may be used to endorse or promote products derived +# from this software without specific prior written permission. +# +# THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY +# EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +# IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR +# PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR +# CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, +# EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, +# PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR +# PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY +# OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +# (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +# OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +# + +# Locate the OptiX distribution. Search relative to the SDK first, then look in the system. + +# Our initial guess will be within the SDK. + +if (WIN32) + set(OptiX_INSTALL_DIR "C:/ProgramData/NVIDIA Corporation/OptiX SDK 5.1.0" CACHE PATH "Path to OptiX installed location.") +else() + set(OptiX_INSTALL_DIR $ENV{OptiX_INSTALL_DIR} CACHE PATH "Path to OptiX installed location.") +endif() + +# The distribution contains both 32 and 64 bit libraries. Adjust the library +# search path based on the bit-ness of the build. (i.e. 64: bin64, lib64; 32: +# bin, lib). Note that on Mac, the OptiX library is a universal binary, so we +# only need to look in lib and not lib64 for 64 bit builds. +if(CMAKE_SIZEOF_VOID_P EQUAL 8 AND NOT APPLE) + set(bit_dest "64") +else() + set(bit_dest "") +endif() + +macro(OPTIX_find_api_library name version) + find_library(${name}_LIBRARY + NAMES ${name}.${version} ${name} + PATHS "${OptiX_INSTALL_DIR}/lib${bit_dest}" + NO_DEFAULT_PATH + ) + find_library(${name}_LIBRARY + NAMES ${name}.${version} ${name} + ) + if(WIN32) + find_file(${name}_DLL + NAMES ${name}.${version}.dll + PATHS "${OptiX_INSTALL_DIR}/bin${bit_dest}" + NO_DEFAULT_PATH + ) + find_file(${name}_DLL + NAMES ${name}.${version}.dll + ) + endif() +endmacro() + +OPTIX_find_api_library(optix 51) +OPTIX_find_api_library(optixu 1) +OPTIX_find_api_library(optix_prime 1) + +# Include +find_path(OptiX_INCLUDE + NAMES optix.h + PATHS "${OptiX_INSTALL_DIR}/include" + NO_DEFAULT_PATH + ) +find_path(OptiX_INCLUDE + NAMES optix.h + ) + +# Check to make sure we found what we were looking for +function(OptiX_report_error error_message required) + if(OptiX_FIND_REQUIRED AND required) + message(FATAL_ERROR "${error_message}") + else() + if(NOT OptiX_FIND_QUIETLY) + message(STATUS "${error_message}") + endif(NOT OptiX_FIND_QUIETLY) + endif() +endfunction() + +if(NOT optix_LIBRARY) + OptiX_report_error("optix library not found. Please locate before proceeding." TRUE) +endif() +if(NOT OptiX_INCLUDE) + OptiX_report_error("OptiX headers (optix.h and friends) not found. Please locate before proceeding." TRUE) +endif() +if(NOT optix_prime_LIBRARY) + OptiX_report_error("optix Prime library not found. Please locate before proceeding." FALSE) +endif() + +# Macro for setting up dummy targets +function(OptiX_add_imported_library name lib_location dll_lib dependent_libs) + set(CMAKE_IMPORT_FILE_VERSION 1) + + # Create imported target + add_library(${name} SHARED IMPORTED) + + # Import target "optix" for configuration "Debug" + if(WIN32) + set_target_properties(${name} PROPERTIES + IMPORTED_IMPLIB "${lib_location}" + #IMPORTED_LINK_INTERFACE_LIBRARIES "glu32;opengl32" + IMPORTED_LOCATION "${dll_lib}" + IMPORTED_LINK_INTERFACE_LIBRARIES "${dependent_libs}" + ) + elseif(UNIX) + set_target_properties(${name} PROPERTIES + #IMPORTED_LINK_INTERFACE_LIBRARIES "glu32;opengl32" + IMPORTED_LOCATION "${lib_location}" + # We don't have versioned filenames for now, and it may not even matter. + #IMPORTED_SONAME "${optix_soname}" + IMPORTED_LINK_INTERFACE_LIBRARIES "${dependent_libs}" + ) + else() + # Unknown system, but at least try and provide the minimum required + # information. + set_target_properties(${name} PROPERTIES + IMPORTED_LOCATION "${lib_location}" + IMPORTED_LINK_INTERFACE_LIBRARIES "${dependent_libs}" + ) + endif() + + # Commands beyond this point should not need to know the version. + set(CMAKE_IMPORT_FILE_VERSION) +endfunction() + +# Sets up a dummy target +OptiX_add_imported_library(optix "${optix_LIBRARY}" "${optix_DLL}" "${OPENGL_LIBRARIES}") +OptiX_add_imported_library(optixu "${optixu_LIBRARY}" "${optixu_DLL}" "") +OptiX_add_imported_library(optix_prime "${optix_prime_LIBRARY}" "${optix_prime_DLL}" "") + +macro(OptiX_check_same_path libA libB) + if(_optix_path_to_${libA}) + if(NOT _optix_path_to_${libA} STREQUAL _optix_path_to_${libB}) + # ${libA} and ${libB} are in different paths. Make sure there isn't a ${libA} next + # to the ${libB}. + get_filename_component(_optix_name_of_${libA} "${${libA}_LIBRARY}" NAME) + if(EXISTS "${_optix_path_to_${libB}}/${_optix_name_of_${libA}}") + message(WARNING " ${libA} library found next to ${libB} library that is not being used. Due to the way we are using rpath, the copy of ${libA} next to ${libB} will be used during loading instead of the one you intended. Consider putting the libraries in the same directory or moving ${_optix_path_to_${libB}}/${_optix_name_of_${libA} out of the way.") + endif() + endif() + set( _${libA}_rpath "-Wl,-rpath,${_optix_path_to_${libA}}" ) + endif() +endmacro() + +# Since liboptix.1.dylib is built with an install name of @rpath, we need to +# compile our samples with the rpath set to where optix exists. +if(APPLE) + get_filename_component(_optix_path_to_optix "${optix_LIBRARY}" PATH) + if(_optix_path_to_optix) + set( _optix_rpath "-Wl,-rpath,${_optix_path_to_optix}" ) + endif() + get_filename_component(_optix_path_to_optixu "${optixu_LIBRARY}" PATH) + OptiX_check_same_path(optixu optix) + get_filename_component(_optix_path_to_optix_prime "${optix_prime_LIBRARY}" PATH) + OptiX_check_same_path(optix_prime optix) + OptiX_check_same_path(optix_prime optixu) + + set( optix_rpath ${_optix_rpath} ${_optixu_rpath} ${_optix_prime_rpath} ) + list(REMOVE_DUPLICATES optix_rpath) +endif() + diff --git a/cmake/configure_optix.cmake b/cmake/configure_optix.cmake new file mode 100644 index 0000000..a2de2b0 --- /dev/null +++ b/cmake/configure_optix.cmake @@ -0,0 +1,68 @@ +# ======================================================================== # +# Copyright 2018 Ingo Wald # +# # +# Licensed under the Apache License, Version 2.0 (the "License"); # +# you may not use this file except in compliance with the License. # +# You may obtain a copy of the License at # +# # +# http://www.apache.org/licenses/LICENSE-2.0 # +# # +# Unless required by applicable law or agreed to in writing, software # +# distributed under the License is distributed on an "AS IS" BASIS, # +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. # +# See the License for the specific language governing permissions and # +# limitations under the License. # +# ======================================================================== # + +set(CMAKE_MODULE_PATH + "${PROJECT_SOURCE_DIR}/cmake" +# "${CMAKE_CURRENT_SOURCE_DIR}/../cmake" + ${CMAKE_MODULE_PATH} + ) + +find_package(CUDA REQUIRED) +find_package(OptiX REQUIRED) + +#include_directories(${CUDA_TOOLKIT_INCLUDE}) +if (CUDA_TOOLKIT_ROOT_DIR) + include_directories(${CUDA_TOOLKIT_ROOT_DIR}/include) +endif() +include_directories(${OptiX_INCLUDE}) + +if (WIN32) + add_definitions(-DNOMINMAX) +endif() + +find_program(BIN2C bin2c + DOC "Path to the cuda-sdk bin2c executable.") + +# this macro defines cmake rules that execute the following four steps: +# 1) compile the given cuda file ${cuda_file} to an intermediary PTX file +# 2) use the 'bin2c' tool (that comes with CUDA) to +# create a second intermediary (.c-)file which defines a const string variable +# (named '${c_var_name}') whose (constant) value is the PTX output +# from the previous step. +# 3) compile the given .c file to an intermediary object file (why thus has +# that PTX string 'embedded' as a global constant. +# 4) assign the name of the intermediary .o file to the cmake variable +# 'output_var', which can then be added to cmake targets. +macro(cuda_compile_and_embed output_var cuda_file) + set(c_var_name ${output_var}) + cuda_compile_ptx(ptx_files ${cuda_file}) + list(GET ptx_files 0 ptx_file) + set(embedded_file ${ptx_file}_embedded.c) +# message("adding rule to compile and embed ${cuda_file} to \"const char ${var_name}[];\"") + add_custom_command( + OUTPUT ${embedded_file} + COMMAND ${BIN2C} -c --padd 0 --type char --name ${c_var_name} ${ptx_file} > ${embedded_file} + DEPENDS ${ptx_file} + COMMENT "compiling (and embedding ptx from) ${cuda_file}" + ) + set(${output_var} ${embedded_file}) +endmacro() + +include_directories(${OptiX_INCLUDE}) + +add_definitions(-D__CUDA_INCLUDE_COMPILER_INTERNAL_HEADERS__=1) + +