initial import

.gitignore

@@ -0,0 +1,2 @@
+*~
+

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)

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}
+  )
+

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;
+};

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;
+};