From ee32c11b396ad931fc4bb275c18407835c9862c1 Mon Sep 17 00:00:00 2001 From: Luc Forget Date: Tue, 13 Apr 2021 09:09:13 -0700 Subject: [PATCH 01/14] Replace float with real_t --- include/camera.hpp | 2 +- include/constant_medium.hpp | 2 +- include/hitable.hpp | 6 ++--- include/material.hpp | 12 +++++----- include/ray.hpp | 2 +- include/render.hpp | 4 ++-- include/rtweekend.hpp | 44 ++++++++++++++++++++----------------- include/sphere.hpp | 2 +- include/texture.hpp | 12 +++++----- include/vec.hpp | 6 ++--- src/main.cpp | 8 +++---- 11 files changed, 51 insertions(+), 49 deletions(-) diff --git a/include/camera.hpp b/include/camera.hpp index cc15f7f..8265bcf 100644 --- a/include/camera.hpp +++ b/include/camera.hpp @@ -96,7 +96,7 @@ class camera { return { origin + offset, lower_left_corner + s * horizontal + t * vertical - origin - offset, - rng.float_t(time0, time1) }; + rng.real(time0, time1) }; } }; diff --git a/include/constant_medium.hpp b/include/constant_medium.hpp index c4b26e8..af195c8 100644 --- a/include/constant_medium.hpp +++ b/include/constant_medium.hpp @@ -62,7 +62,7 @@ class constant_medium { /// Distance between the two hitpoints affect of probability /// of the ray hitting a smoke particle const auto distance_inside_boundary = (rec2.t - rec1.t) * ray_length; - const auto hit_distance = neg_inv_density * sycl::log(rng.float_t()); + const auto hit_distance = neg_inv_density * sycl::log(rng.real()); /// With lower density, hit_distance has higher probabilty /// of being greater than distance_inside_boundary diff --git a/include/hitable.hpp b/include/hitable.hpp index d5eb79a..76109d0 100644 --- a/include/hitable.hpp +++ b/include/hitable.hpp @@ -7,14 +7,14 @@ class hit_record { public: - float t; // + real_t t; // point p; // hit point vec normal; // normal at hit point bool front_face; // to check if hit point is on the outer surface /*local coordinates for rectangles and mercator coordintes for spheres */ - float u; - float v; + real_t u; + real_t v; // To set if the hit point is on the front face void set_face_normal(const ray& r, const vec& outward_normal) { diff --git a/include/material.hpp b/include/material.hpp index d11bffb..c3c0e7d 100644 --- a/include/material.hpp +++ b/include/material.hpp @@ -32,7 +32,7 @@ struct lambertian_material { struct metal_material { metal_material() = default; - metal_material(const color& a, float f) + metal_material(const color& a, real_t f) : albedo { a } , fuzz { std::clamp(f, 0.0f, 1.0f) } {} @@ -49,7 +49,7 @@ struct metal_material { color emitted(auto&, const hit_record& rec) { return color(0, 0, 0); } color albedo; - float fuzz; + real_t fuzz; }; struct dielectric_material { @@ -71,14 +71,14 @@ struct dielectric_material { // at hit point auto& rng = ctx.rng; attenuation *= albedo; - float refraction_ratio = rec.front_face ? (1.0f / ref_idx) : ref_idx; + real_t refraction_ratio = rec.front_face ? (1.0f / ref_idx) : ref_idx; vec unit_direction = unit_vector(r_in.direction()); - float cos_theta = sycl::fmin(-sycl::dot(unit_direction, rec.normal), 1.0f); - float sin_theta = sycl::sqrt(1.0f - cos_theta * cos_theta); + real_t cos_theta = sycl::fmin(-sycl::dot(unit_direction, rec.normal), 1.0f); + real_t sin_theta = sycl::sqrt(1.0f - cos_theta * cos_theta); bool cannot_refract = refraction_ratio * sin_theta > 1.0f; vec direction; if (cannot_refract || - reflectance(cos_theta, refraction_ratio) > rng.float_t()) + reflectance(cos_theta, refraction_ratio) > rng.real()) direction = reflect(unit_direction, rec.normal); else direction = refract(unit_direction, rec.normal, refraction_ratio); diff --git a/include/ray.hpp b/include/ray.hpp index 902e615..d7ef8fb 100644 --- a/include/ray.hpp +++ b/include/ray.hpp @@ -18,7 +18,7 @@ class ray { // returns point along the ray at distance t from ray's origin // the ray P(t) = Origin + t*direction - point at(float t) const { return orig + t * dir; } + point at(real_t t) const { return orig + t * dir; } public: // To store the origin and direction of the ray diff --git a/include/render.hpp b/include/render.hpp index 565f037..7f8f59a 100644 --- a/include/render.hpp +++ b/include/render.hpp @@ -93,8 +93,8 @@ inline auto render_pixel(auto& ctx, int x_coord, int y_coord, camera const& cam, color final_color(0.0f, 0.0f, 0.0f); for (auto i = 0; i < samples; i++) { - const auto u = (x_coord + rng.float_t()) / width; - const auto v = (y_coord + rng.float_t()) / height; + const auto u = (x_coord + rng.real()) / width; + const auto v = (y_coord + rng.real()) / height; // u and v are points on the viewport ray r = cam.get_ray(u, v, rng); final_color += get_color(r); diff --git a/include/rtweekend.hpp b/include/rtweekend.hpp index db7618d..cf0871d 100644 --- a/include/rtweekend.hpp +++ b/include/rtweekend.hpp @@ -16,19 +16,23 @@ #include #include +using real_t = float; + // Constants -constexpr float infinity = std::numeric_limits::infinity(); -constexpr float pi = 3.1415926535897932385f; +constexpr real_t infinity = std::numeric_limits::infinity(); +constexpr real_t pi = 3.1415926535897932385f; // type aliases for float3 - vec, point and color -using point = sycl::float3; -using color = sycl::float3; -using vec = sycl::float3; +using real_vec = sycl::float3; + +using point = real_vec; +using color = real_vec; +using vec = real_vec; // Utility Functions -inline float degrees_to_radians(float degrees) { return degrees * pi / 180.0f; } +inline real_t degrees_to_radians(real_t degrees) { return degrees * pi / 180.0f; } class LocalPseudoRNG { public: @@ -36,42 +40,42 @@ class LocalPseudoRNG { : generator { init_state } {} // Returns a random float in 0.f 1. - inline float float_t() { - constexpr float scale = 1.f / (uint64_t { 1 } << 32); + inline real_t real() { + constexpr real_t scale = 1.f / (uint64_t { 1 } << 32); return generator() * scale; } // Returns a random float in min, max - inline float float_t(float min, float max) { + inline real_t real(real_t min, real_t max) { // TODO use FMA ? - return min + (max - min) * float_t(); + return min + (max - min) * real(); } // Returns a random vector with coordinates in 0.f 1. - inline vec vec_t() { return { float_t(), float_t(), float_t() }; } + inline vec vec_t() { return { real(), real(), real() }; } // Returns a random vec with coordinates in min, max - inline vec vec_t(float min, float max) { + inline vec vec_t(real_t min, real_t max) { auto scale = max - min; return vec_t() * scale + min; } // Returns a random unit vector inline vec unit_vec() { - auto x = float_t(-1.f, 1.f); + auto x = real(-1.f, 1.f); auto maxy = sycl::sqrt(1 - x * x); - auto y = float_t(-maxy, maxy); + auto y = real(-maxy, maxy); auto absz = sycl::sqrt(maxy * maxy - y * y); - auto z = (float_t() > 0.5) ? absz : -absz; + auto z = (real() > 0.5) ? absz : -absz; return vec(x, y, z); } // Returns a random vector in the unit ball of usual norm inline vec in_unit_ball() { // Polar coordinates r, theta, phi - auto r = float_t(); - auto theta = float_t(0, 2 * pi); - auto phi = float_t(0, pi); + auto r = real(); + auto theta = real(0, 2 * pi); + auto phi = real(0, pi); auto plan_seed = r * sycl::sin(phi); auto z = r * sycl::cos(phi); @@ -81,9 +85,9 @@ class LocalPseudoRNG { // Return a random vector in the unit disk of usual norm in plane x, y inline vec in_unit_disk() { - auto x = float_t(-1.f, 1.f); + auto x = real(-1.f, 1.f); auto maxy = sycl::sqrt(1 - x * x); - auto y = float_t(-maxy, maxy); + auto y = real(-maxy, maxy); return { x, y, 0.f }; } diff --git a/include/sphere.hpp b/include/sphere.hpp index ccc5b26..33624f1 100644 --- a/include/sphere.hpp +++ b/include/sphere.hpp @@ -10,7 +10,7 @@ /* Computes normalised values of theta and phi. The input vector p corresponds to a vector passing through the centre of the a sphere and the hipoint on the surface of the sphere */ -std::pair mercator_coordinates(const vec& p) { +std::pair mercator_coordinates(const vec& p) { // phi is the angle around the axis auto phi = sycl::atan2(p.z(), p.x()); // theta is the angle down from the pole diff --git a/include/texture.hpp b/include/texture.hpp index 50002cc..db1c145 100644 --- a/include/texture.hpp +++ b/include/texture.hpp @@ -19,7 +19,7 @@ struct solid_texture { solid_texture() = default; solid_texture(const color& c) : color_value { c } {} - solid_texture(float red, float green, float blue) + solid_texture(real_t red, real_t green, real_t blue) : solid_texture { color { red, green, blue } } {} // For solid texture, the color is same throughout the sphere color value(auto&, const hit_record&) const { return color_value; } @@ -77,10 +77,10 @@ struct image_texture { std::size_t offset; /// The repetition rate of the image - float cyclic_frequency { 1.f }; + real_t cyclic_frequency { 1.f }; image_texture(std::size_t _width, std::size_t _height, std::size_t _offset, - float _cyclic_frequency) + real_t _cyclic_frequency) : width { _width } , height { _height } , offset { _offset } @@ -95,7 +95,7 @@ struct image_texture { the image in the texture */ static image_texture image_texture_factory(const char* file_name, - float _cyclic_frequency = 1) { + real_t _cyclic_frequency = 1) { assert(!frozen); auto components_per_pixel = bytes_per_pixel; uint8_t* _data; @@ -137,10 +137,10 @@ struct image_texture { // The image is repeated by the repetition factor std::size_t i = - sycl::fmod(rec.u * cyclic_frequency, (float)1) * (width - 1); + sycl::fmod(rec.u * cyclic_frequency, (real_t)1) * (width - 1); // The image frame buffer is going downwards, so flip the y axis std::size_t j = - (1 - sycl::fmod(rec.v * cyclic_frequency, (float)1)) * (height - 1); + (1 - sycl::fmod(rec.v * cyclic_frequency, (real_t)1)) * (height - 1); std::size_t local_offset = j * width + i; std::size_t pix_idx = local_offset + offset; auto scale = 1.f / 255; diff --git a/include/vec.hpp b/include/vec.hpp index f9327ac..98c7d5d 100644 --- a/include/vec.hpp +++ b/include/vec.hpp @@ -5,10 +5,8 @@ #include #include -using real_t = float; - // vec Utility Functions -inline float length_squared(const vec& v) { +inline real_t length_squared(const vec& v) { return sycl::fma(v.x(), v.x(), sycl::fma(v.y(), v.y(), v.z() * v.z())); } @@ -26,7 +24,7 @@ inline vec unit_vector(const vec& v) { return v / sycl::length(v); } vec reflect(const vec& v, const vec& n) { return v - 2 * sycl::dot(v, n) * n; } // Computes refracted ray's direction based on refractive index -vec refract(const vec& uv, const vec& n, float etai_over_etat) { +vec refract(const vec& uv, const vec& n, real_t etai_over_etat) { auto cos_theta = sycl::fmin(-sycl::dot(uv, n), 1.0f); vec r_out_perp = etai_over_etat * (uv + cos_theta * n); vec r_out_parallel = diff --git a/src/main.cpp b/src/main.cpp index df7d03a..2456c25 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -78,9 +78,9 @@ int main() { for (int a = -11; a < 11; a++) { for (int b = -11; b < 11; b++) { // Based on a random variable , the material type is chosen - auto choose_mat = rng.float_t(); + auto choose_mat = rng.real(); // Spheres are placed at a point randomly displaced from a,b - point center(a + 0.9f * rng.float_t(), 0.2f, b + 0.9f * rng.float_t()); + point center(a + 0.9f * rng.real(), 0.2f, b + 0.9f * rng.real()); if (sycl::length((center - point(4, 0.2f, 0))) > 0.9f) { if (choose_mat < 0.4f) { // Lambertian @@ -90,13 +90,13 @@ int main() { } else if (choose_mat < 0.8f) { // Lambertian movig spheres auto albedo = rng.vec_t() * rng.vec_t(); - auto center2 = center + point { 0, rng.float_t(0, 0.25f), 0 }; + auto center2 = center + point { 0, rng.real(0, 0.25f), 0 }; hittables.emplace_back(sphere(center, center2, 0.0f, 1.0f, 0.2f, lambertian_material(albedo))); } else if (choose_mat < 0.95f) { // metal auto albedo = rng.vec_t(0.5f, 1); - auto fuzz = rng.float_t(0, 0.5f); + auto fuzz = rng.real(0, 0.5f); hittables.emplace_back( sphere(center, 0.2f, metal_material(albedo, fuzz))); } else { From 3979043d2556cf2dc3632d7ca1af87eacf876e8d Mon Sep 17 00:00:00 2001 From: Luc Forget Date: Wed, 14 Apr 2021 06:08:05 -0700 Subject: [PATCH 02/14] Adding monostate in front of each variant --- include/box.hpp | 9 ++---- include/constant_medium.hpp | 18 ++++------- include/hit_record.hpp | 23 +++++++++++++ include/hitable.hpp | 53 +++++++++++++++++++++--------- include/material.hpp | 64 +++++++++++++++++++++++++++++++------ include/render.hpp | 29 ++++------------- include/texture.hpp | 22 +++++++++++-- 7 files changed, 151 insertions(+), 67 deletions(-) create mode 100644 include/hit_record.hpp diff --git a/include/box.hpp b/include/box.hpp index feea59f..65042a4 100644 --- a/include/box.hpp +++ b/include/box.hpp @@ -34,12 +34,9 @@ class box { auto closest_so_far = max; // Checking if the ray hits any of the sides for (const auto& side : sides) { - if (dev_visit( - [&](auto&& arg) { - return arg.hit(ctx, r, min, closest_so_far, temp_rec, - temp_material_type); - }, - side)) { + if (dev_visit(hittable_hit_visitor(ctx, r, min, closest_so_far, temp_rec, + temp_material_type), + side)) { hit_anything = true; closest_so_far = temp_rec.t; rec = temp_rec; diff --git a/include/constant_medium.hpp b/include/constant_medium.hpp index af195c8..0408698 100644 --- a/include/constant_medium.hpp +++ b/include/constant_medium.hpp @@ -7,7 +7,7 @@ #include "texture.hpp" #include "visit.hpp" -using hittableVolume_t = std::variant; +using hittableVolume_t = std::variant; /** * A ray going through the volume can either make it all the way through @@ -31,20 +31,14 @@ class constant_medium { hit_material_type = phase_function; material_t temp_material_type; hit_record rec1, rec2; - if (!dev_visit( - [&](auto&& arg) { - return arg.hit(ctx, r, -infinity, infinity, rec1, - temp_material_type); - }, - boundary)) { + if (!dev_visit(hittable_hit_visitor(ctx, r, -infinity, infinity, rec1, + temp_material_type), + boundary)) { return false; } - if (!dev_visit( - [&](auto&& arg) { - return arg.hit(ctx, r, rec1.t + 0.0001f, infinity, rec2, - temp_material_type); - }, + if (!dev_visit(hittable_hit_visitor(ctx, r, rec1.t + 0.0001f, infinity, rec2, + temp_material_type), boundary)) { return false; } diff --git a/include/hit_record.hpp b/include/hit_record.hpp new file mode 100644 index 0000000..342fb81 --- /dev/null +++ b/include/hit_record.hpp @@ -0,0 +1,23 @@ +#ifndef HIT_RECORD_HPP +#define HIT_RECORD_HPP + +#include "rtweekend.hpp" + +class hit_record { + public: + real_t t; // + point p; // hit point + vec normal; // normal at hit point + bool front_face; // to check if hit point is on the outer surface + /*local coordinates for rectangles + and mercator coordintes for spheres */ + real_t u; + real_t v; + + // To set if the hit point is on the front face + void set_face_normal(const ray& r, const vec& outward_normal) { + front_face = dot(r.direction(), outward_normal) < 0; + normal = front_face ? outward_normal : vec {} - outward_normal; + } +}; +#endif \ No newline at end of file diff --git a/include/hitable.hpp b/include/hitable.hpp index 76109d0..d12c15c 100644 --- a/include/hitable.hpp +++ b/include/hitable.hpp @@ -1,26 +1,49 @@ #ifndef HITTABLE_H #define HITTABLE_H -#include "ray.hpp" +#include + +#include "hit_record.hpp" +#include "material.hpp" #include "rtweekend.hpp" #include "vec.hpp" -class hit_record { +struct hittable_hit_visitor { + private: + task_context& ctx; + const ray& r; + real_t min; + real_t max; + hit_record& rec; + material_t& hit_material_type; + public: - real_t t; // - point p; // hit point - vec normal; // normal at hit point - bool front_face; // to check if hit point is on the outer surface - /*local coordinates for rectangles - and mercator coordintes for spheres */ - real_t u; - real_t v; - - // To set if the hit point is on the front face - void set_face_normal(const ray& r, const vec& outward_normal) { - front_face = dot(r.direction(), outward_normal) < 0; - normal = front_face ? outward_normal : vec {} - outward_normal; + hittable_hit_visitor(task_context& ctx, const ray& r, real_t min, real_t max, + hit_record& rec, material_t& hit_material_type) + : ctx { ctx } + , r { r } + , min { min } + , max { max } + , rec { rec } + , hit_material_type { hit_material_type } {} + + template bool operator()(H&& hittable) { + return hittable.hit(ctx, r, min, max, rec, hit_material_type); + } + + bool operator()(std::monostate) { + assert(fase && "unreachable"); + return false; } }; +#include "box.hpp" +#include "constant_medium.hpp" +#include "ray.hpp" +#include "rectangle.hpp" +#include "sphere.hpp" +#include "triangle.hpp" + +using hittable_t = std::variant; #endif diff --git a/include/material.hpp b/include/material.hpp index c3c0e7d..9622a4f 100644 --- a/include/material.hpp +++ b/include/material.hpp @@ -3,7 +3,7 @@ #include -#include "hitable.hpp" +#include "hit_record.hpp" #include "texture.hpp" #include "vec.hpp" #include "visit.hpp" @@ -22,8 +22,7 @@ struct lambertian_material { scattered = ray(rec.p, scatter_direction, r_in.time()); // Attenuation of the ray hitting the object is modified based on the color // at hit point - attenuation *= - dev_visit([&](auto&& arg) { return arg.value(ctx, rec); }, albedo); + attenuation *= dev_visit(texture_value_visitor(ctx, rec), albedo); return true; } color emitted(auto&, const hit_record& rec) { return color(0, 0, 0); } @@ -77,8 +76,7 @@ struct dielectric_material { real_t sin_theta = sycl::sqrt(1.0f - cos_theta * cos_theta); bool cannot_refract = refraction_ratio * sin_theta > 1.0f; vec direction; - if (cannot_refract || - reflectance(cos_theta, refraction_ratio) > rng.real()) + if (cannot_refract || reflectance(cos_theta, refraction_ratio) > rng.real()) direction = reflect(unit_direction, rec.normal); else direction = refract(unit_direction, rec.normal, refraction_ratio); @@ -104,7 +102,7 @@ struct lightsource_material { template bool scatter(T&...) const { return false; } color emitted(auto& ctx, const hit_record& rec) { - return dev_visit([&](auto&& arg) { return arg.value(ctx, rec); }, emit); + return dev_visit(texture_value_visitor(ctx, rec), emit); } texture_t emit; @@ -120,8 +118,7 @@ struct isotropic_material { color& attenuation, ray& scattered) const { auto& rng = ctx.rng; scattered = ray(rec.p, rng.in_unit_ball(), r_in.time()); - attenuation *= - dev_visit([&](auto&& arg) { return arg.value(ctx, rec); }, albedo); + attenuation *= dev_visit(texture_value_visitor(ctx, rec), albedo); return true; } @@ -131,7 +128,54 @@ struct isotropic_material { }; using material_t = - std::variant; + std::variant; + +struct material_emitted_visitor { + private: + task_context& ctx; + const hit_record& rec; + + public: + material_emitted_visitor(task_context& ctx, hit_record& rec) + : ctx { ctx } + , rec { rec } {} + + template color operator()(M&& material) { + return material.emitted(ctx, rec); + } + + color operator()(std::monostate) { + assert(false && "unreachable"); + return { 0.f, 0.f, 0.f }; + } +}; + +struct material_scatter_visitor { + private: + task_context& ctx; + const ray& r_in; + const hit_record& rec; + color& attenuation; + ray& scattered; + + public: + material_scatter_visitor(auto& ctx, const ray& r_in, const hit_record& rec, + color& attenuation, ray& scattered) + : ctx { ctx } + , r_in { r_in } + , rec { rec } + , attenuation { attenuation } + , scattered { scattered } {} + + template bool operator()(M&& material) { + return material.scatter(ctx, r_in, rec, attenuation, scattered); + } + + bool operator()(std::monostate) { + assert(false && "unreachable"); + return false; + } +}; #endif diff --git a/include/render.hpp b/include/render.hpp index 7f8f59a..b6dd4c4 100644 --- a/include/render.hpp +++ b/include/render.hpp @@ -3,25 +3,17 @@ #include #include -#include "box.hpp" #include "build_parameters.hpp" #include "camera.hpp" -#include "constant_medium.hpp" #include "hitable.hpp" #include "material.hpp" #include "ray.hpp" -#include "rectangle.hpp" #include "rtweekend.hpp" -#include "sphere.hpp" #include "sycl.hpp" #include "texture.hpp" -#include "triangle.hpp" #include "vec.hpp" #include "visit.hpp" -using hittable_t = - std::variant; - template inline auto render_pixel(auto& ctx, int x_coord, int y_coord, camera const& cam, auto& hittable_acc, auto fb_acc) { @@ -35,12 +27,9 @@ inline auto render_pixel(auto& ctx, int x_coord, int y_coord, camera const& cam, auto closest_so_far = infinity; // Checking if the ray hits any of the spheres for (auto i = 0; i < hittable_acc.get_count(); i++) { - if (dev_visit( - [&](auto&& arg) { - return arg.hit(ctx, r, 0.001f, closest_so_far, temp_rec, - temp_material_type); - }, - hittable_acc[i])) { + if (dev_visit(hittable_hit_visitor(ctx, r, 0.001f, closest_so_far, + temp_rec, temp_material_type), + hittable_acc[i])) { hit_anything = true; closest_so_far = temp_rec.t; rec = temp_rec; @@ -58,14 +47,10 @@ inline auto render_pixel(auto& ctx, int x_coord, int y_coord, camera const& cam, for (auto i = 0; i < depth; i++) { hit_record rec; if (hit_world(cur_ray, rec, material_type)) { - emitted = dev_visit([&](auto&& arg) { return arg.emitted(ctx, rec); }, - material_type); - if (dev_visit( - [&](auto&& arg) { - return arg.scatter(ctx, cur_ray, rec, cur_attenuation, - scattered); - }, - material_type)) { + emitted = dev_visit(material_emitted_visitor(ctx, rec), material_type); + if (dev_visit(material_scatter_visitor(ctx, cur_ray, rec, + cur_attenuation, scattered), + material_type)) { // On hitting the object, the ray gets scattered cur_ray = scattered; } else { diff --git a/include/texture.hpp b/include/texture.hpp index db1c145..c547f57 100644 --- a/include/texture.hpp +++ b/include/texture.hpp @@ -1,6 +1,6 @@ #ifndef RT_SYCL_TEXTURE_HPP #define RT_SYCL_TEXTURE_HPP -#include "hitable.hpp" +#include "hit_record.hpp" #include "rtweekend.hpp" #include "vec.hpp" #include @@ -151,7 +151,25 @@ struct image_texture { } }; -using texture_t = std::variant; +using texture_t = std::variant; + +struct texture_value_visitor { + private: + task_context& ctx; + const hit_record& rec; + + public: + texture_value_visitor(task_context& ctx, const hit_record& rec):ctx{ctx}, rec{rec}{} + template + color operator()(T&& texture) { + return texture.value(ctx, rec); + } + + color operator()(std::monostate) { + assert(false && "unreachable"); + return {0.f,0.f,0.f}; + } +}; // Start filled with the fallback texture (solid blue) for texture load error std::vector image_texture::texture_data { 0, 0, 1 }; From 1893a478812812bc37e45613330ac23417a242a9 Mon Sep 17 00:00:00 2001 From: Luc Forget Date: Tue, 20 Apr 2021 04:24:50 -0700 Subject: [PATCH 03/14] Replace visitor object by function --- include/box.hpp | 8 +++-- include/constant_medium.hpp | 18 +++++++--- include/hitable.hpp | 36 -------------------- include/material.hpp | 67 ++++++++----------------------------- include/render.hpp | 22 +++++++++--- include/texture.hpp | 18 ---------- include/visit.hpp | 11 ++++++ 7 files changed, 61 insertions(+), 119 deletions(-) diff --git a/include/box.hpp b/include/box.hpp index 65042a4..cfa80bd 100644 --- a/include/box.hpp +++ b/include/box.hpp @@ -34,8 +34,12 @@ class box { auto closest_so_far = max; // Checking if the ray hits any of the sides for (const auto& side : sides) { - if (dev_visit(hittable_hit_visitor(ctx, r, min, closest_so_far, temp_rec, - temp_material_type), + if (dev_visit(monostate_dispatch( + [&](auto&& object) { + return object.hit(ctx, r, min, closest_so_far, + temp_rec, temp_material_type); + }, + false), side)) { hit_anything = true; closest_so_far = temp_rec.t; diff --git a/include/constant_medium.hpp b/include/constant_medium.hpp index 0408698..ebc8fc7 100644 --- a/include/constant_medium.hpp +++ b/include/constant_medium.hpp @@ -31,15 +31,23 @@ class constant_medium { hit_material_type = phase_function; material_t temp_material_type; hit_record rec1, rec2; - if (!dev_visit(hittable_hit_visitor(ctx, r, -infinity, infinity, rec1, - temp_material_type), + if (!dev_visit(monostate_dispatch( + [&](auto&& object) { + return object.hit(ctx, r, -infinity, infinity, rec1, + temp_material_type); + }, + false), boundary)) { return false; } - if (!dev_visit(hittable_hit_visitor(ctx, r, rec1.t + 0.0001f, infinity, rec2, - temp_material_type), - boundary)) { + if (!dev_visit(monostate_dispatch( + [&](auto&& object) { + return object.hit(ctx, r, rec1.t + 0.0001f, infinity, + rec2, temp_material_type); + }, + false), + boundary)) { return false; } diff --git a/include/hitable.hpp b/include/hitable.hpp index d12c15c..dfd8d39 100644 --- a/include/hitable.hpp +++ b/include/hitable.hpp @@ -1,42 +1,6 @@ #ifndef HITTABLE_H #define HITTABLE_H -#include - -#include "hit_record.hpp" -#include "material.hpp" -#include "rtweekend.hpp" -#include "vec.hpp" - -struct hittable_hit_visitor { - private: - task_context& ctx; - const ray& r; - real_t min; - real_t max; - hit_record& rec; - material_t& hit_material_type; - - public: - hittable_hit_visitor(task_context& ctx, const ray& r, real_t min, real_t max, - hit_record& rec, material_t& hit_material_type) - : ctx { ctx } - , r { r } - , min { min } - , max { max } - , rec { rec } - , hit_material_type { hit_material_type } {} - - template bool operator()(H&& hittable) { - return hittable.hit(ctx, r, min, max, rec, hit_material_type); - } - - bool operator()(std::monostate) { - assert(fase && "unreachable"); - return false; - } -}; - #include "box.hpp" #include "constant_medium.hpp" #include "ray.hpp" diff --git a/include/material.hpp b/include/material.hpp index 9622a4f..a999fc6 100644 --- a/include/material.hpp +++ b/include/material.hpp @@ -22,10 +22,13 @@ struct lambertian_material { scattered = ray(rec.p, scatter_direction, r_in.time()); // Attenuation of the ray hitting the object is modified based on the color // at hit point - attenuation *= dev_visit(texture_value_visitor(ctx, rec), albedo); + attenuation *= dev_visit( + monostate_dispatch([&](auto&& t) { return t.value(ctx, rec); }, + color { 0.f, 0.f, 0.f }), + albedo); return true; } - color emitted(auto&, const hit_record& rec) { return color(0, 0, 0); } + color emitted(auto&, const hit_record& rec) { return color(0.f, 0.f, 0.f); } texture_t albedo; }; @@ -46,7 +49,7 @@ struct metal_material { return (dot(scattered.direction(), rec.normal) > 0); } - color emitted(auto&, const hit_record& rec) { return color(0, 0, 0); } + color emitted(auto&, const hit_record& rec) { return color(0.f, 0.f, 0.f); } color albedo; real_t fuzz; }; @@ -102,7 +105,10 @@ struct lightsource_material { template bool scatter(T&...) const { return false; } color emitted(auto& ctx, const hit_record& rec) { - return dev_visit(texture_value_visitor(ctx, rec), emit); + return dev_visit( + monostate_dispatch([&](auto&& t) { return t.value(ctx, rec); }, + color { 0.f, 0.f, 0.f }), + emit); } texture_t emit; @@ -118,7 +124,10 @@ struct isotropic_material { color& attenuation, ray& scattered) const { auto& rng = ctx.rng; scattered = ray(rec.p, rng.in_unit_ball(), r_in.time()); - attenuation *= dev_visit(texture_value_visitor(ctx, rec), albedo); + attenuation *= + dev_visit(monostate_dispatch([&](auto&& t) { return t.value(ctx, rec); }, + color { 0.f, 0.f, 0.f }), + albedo); return true; } @@ -130,52 +139,4 @@ struct isotropic_material { using material_t = std::variant; - -struct material_emitted_visitor { - private: - task_context& ctx; - const hit_record& rec; - - public: - material_emitted_visitor(task_context& ctx, hit_record& rec) - : ctx { ctx } - , rec { rec } {} - - template color operator()(M&& material) { - return material.emitted(ctx, rec); - } - - color operator()(std::monostate) { - assert(false && "unreachable"); - return { 0.f, 0.f, 0.f }; - } -}; - -struct material_scatter_visitor { - private: - task_context& ctx; - const ray& r_in; - const hit_record& rec; - color& attenuation; - ray& scattered; - - public: - material_scatter_visitor(auto& ctx, const ray& r_in, const hit_record& rec, - color& attenuation, ray& scattered) - : ctx { ctx } - , r_in { r_in } - , rec { rec } - , attenuation { attenuation } - , scattered { scattered } {} - - template bool operator()(M&& material) { - return material.scatter(ctx, r_in, rec, attenuation, scattered); - } - - bool operator()(std::monostate) { - assert(false && "unreachable"); - return false; - } -}; - #endif diff --git a/include/render.hpp b/include/render.hpp index b6dd4c4..0713873 100644 --- a/include/render.hpp +++ b/include/render.hpp @@ -27,8 +27,12 @@ inline auto render_pixel(auto& ctx, int x_coord, int y_coord, camera const& cam, auto closest_so_far = infinity; // Checking if the ray hits any of the spheres for (auto i = 0; i < hittable_acc.get_count(); i++) { - if (dev_visit(hittable_hit_visitor(ctx, r, 0.001f, closest_so_far, - temp_rec, temp_material_type), + if (dev_visit(monostate_dispatch( + [&](auto&& object) { + return object.hit(ctx, r, 0.001f, closest_so_far, + temp_rec, temp_material_type); + }, + false), hittable_acc[i])) { hit_anything = true; closest_so_far = temp_rec.t; @@ -47,9 +51,17 @@ inline auto render_pixel(auto& ctx, int x_coord, int y_coord, camera const& cam, for (auto i = 0; i < depth; i++) { hit_record rec; if (hit_world(cur_ray, rec, material_type)) { - emitted = dev_visit(material_emitted_visitor(ctx, rec), material_type); - if (dev_visit(material_scatter_visitor(ctx, cur_ray, rec, - cur_attenuation, scattered), + emitted = + dev_visit(monostate_dispatch( + [&](auto&& mat) { return mat.emitted(ctx, rec); }, + color { 0.f, 0.f, 0.f }), + material_type); + if (dev_visit(monostate_dispatch( + [&](auto&& mat) { + return mat.scatter(ctx, cur_ray, rec, + cur_attenuation, scattered); + }, + false), material_type)) { // On hitting the object, the ray gets scattered cur_ray = scattered; diff --git a/include/texture.hpp b/include/texture.hpp index c547f57..3753f98 100644 --- a/include/texture.hpp +++ b/include/texture.hpp @@ -153,24 +153,6 @@ struct image_texture { using texture_t = std::variant; -struct texture_value_visitor { - private: - task_context& ctx; - const hit_record& rec; - - public: - texture_value_visitor(task_context& ctx, const hit_record& rec):ctx{ctx}, rec{rec}{} - template - color operator()(T&& texture) { - return texture.value(ctx, rec); - } - - color operator()(std::monostate) { - assert(false && "unreachable"); - return {0.f,0.f,0.f}; - } -}; - // Start filled with the fallback texture (solid blue) for texture load error std::vector image_texture::texture_data { 0, 0, 1 }; bool image_texture::frozen = false; diff --git a/include/visit.hpp b/include/visit.hpp index 9aa8e22..23fde8f 100644 --- a/include/visit.hpp +++ b/include/visit.hpp @@ -45,6 +45,17 @@ decltype(auto) visit_single(Func&& f, Var&& var) { } } // namespace detail +auto monostate_dispatch(auto&& dispatch, auto&& monostate_value) { + return [&](auto&& arg) { + if constexpr (std::is_same_v, std::monostate>) { + assert(false && "Try to dispatch to monostate value"); + return monostate_value; + } else { + return dispatch(arg); + } + }; +} + /// dev_visit is std::visit implementation suitable to be used in device code. /// this version of visit doesn't use any function pointer but uses if series /// which will be turned into switch case by the optimizer. From 174149664c7c8027eaadc8387b9eb658ee7acf91 Mon Sep 17 00:00:00 2001 From: Luc Forget Date: Tue, 20 Apr 2021 05:38:38 -0700 Subject: [PATCH 04/14] Simplified version --- CMakeLists.txt | 7 ++ include/box.hpp | 59 ---------------- include/build_parameters.hpp | 1 + include/constant_medium.hpp | 86 ----------------------- include/hitable.hpp | 6 +- include/material.hpp | 65 +---------------- include/rectangle.hpp | 132 ----------------------------------- include/render.hpp | 12 ++-- include/rtweekend.hpp | 2 - include/texture.hpp | 105 +--------------------------- src/main.cpp | 40 +++-------- 11 files changed, 25 insertions(+), 490 deletions(-) delete mode 100644 include/box.hpp delete mode 100644 include/constant_medium.hpp delete mode 100644 include/rectangle.hpp diff --git a/CMakeLists.txt b/CMakeLists.txt index 92d0ae6..3d0cc4f 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -53,6 +53,12 @@ if(NOT OUTPUT_HEIGHT) STRING "Image height in pixel" FORCE) endif() +if(NOT SAMPLES) + message(STATUS "Setting samples to 32 as none was specified.") + set(SAMPLES "32" CACHE + STRING "Number of ray that are casted for each pixel" FORCE) +endif() + set(SYCL_RT_SRC_DIR ${CMAKE_CURRENT_SOURCE_DIR}/src) set(SYCL_RT_INCLUDE_DIR ${CMAKE_CURRENT_SOURCE_DIR}/include) @@ -62,6 +68,7 @@ add_executable(sycl-rt ${SYCL_RT_SRC_DIR}/main.cpp) target_include_directories(sycl-rt PRIVATE ${SYCL_RT_INCLUDE_DIR}) target_compile_definitions(sycl-rt PRIVATE OUTPUT_WIDTH=${OUTPUT_WIDTH}) target_compile_definitions(sycl-rt PRIVATE OUTPUT_HEIGHT=${OUTPUT_HEIGHT}) +target_compile_definitions(sycl-rt PRIVATE SAMPLES=${SAMPLES}) # This is a SYCL program if ("${SYCL_CXX_COMPILER}" STREQUAL "") diff --git a/include/box.hpp b/include/box.hpp deleted file mode 100644 index cfa80bd..0000000 --- a/include/box.hpp +++ /dev/null @@ -1,59 +0,0 @@ -#ifndef BOX_HPP -#define BOX_HPP - -#include "rectangle.hpp" -#include "rtweekend.hpp" -#include "visit.hpp" - -/// This class implements a axis aligned cuboid using 6 rectangles -class box { - public: - box() = default; - - /// p0 = { x0, y0, z0 } and p1 = { x1, y1. z1 } - /// where x0 <= x1, y0 <= y1 and z0 <= z1 - box(const point& p0, const point& p1, const material_t& mat_type) - : box_min { p0 } - , box_max { p1 } - , material_type { mat_type } { - /// Add six sides of the box based on box_min and box_max to sides - sides[0] = xy_rect(p0.x(), p1.x(), p0.y(), p1.y(), p1.z(), mat_type); - sides[1] = xy_rect(p0.x(), p1.x(), p0.y(), p1.y(), p0.z(), mat_type); - sides[2] = xz_rect(p0.x(), p1.x(), p0.z(), p1.z(), p1.y(), mat_type); - sides[3] = xz_rect(p0.x(), p1.x(), p0.z(), p1.z(), p0.y(), mat_type); - sides[4] = yz_rect(p0.y(), p1.y(), p0.z(), p1.z(), p1.x(), mat_type); - sides[5] = yz_rect(p0.y(), p1.y(), p0.z(), p1.z(), p0.x(), mat_type); - } - - /// Compute ray interaction with the box - bool hit(auto& ctx, const ray& r, real_t min, real_t max, hit_record& rec, - material_t& hit_material_type) const { - hit_record temp_rec; - material_t temp_material_type; - auto hit_anything = false; - auto closest_so_far = max; - // Checking if the ray hits any of the sides - for (const auto& side : sides) { - if (dev_visit(monostate_dispatch( - [&](auto&& object) { - return object.hit(ctx, r, min, closest_so_far, - temp_rec, temp_material_type); - }, - false), - side)) { - hit_anything = true; - closest_so_far = temp_rec.t; - rec = temp_rec; - hit_material_type = temp_material_type; - } - } - return hit_anything; - } - - point box_min; - point box_max; - material_t material_type; - std::array sides; -}; - -#endif diff --git a/include/build_parameters.hpp b/include/build_parameters.hpp index 5ecc388..84d8b03 100644 --- a/include/build_parameters.hpp +++ b/include/build_parameters.hpp @@ -16,6 +16,7 @@ constexpr bool use_sycl_compiler = false; constexpr int output_width = OUTPUT_WIDTH; constexpr int output_height = OUTPUT_HEIGHT; +constexpr int samples = SAMPLES; } // namespace buildparams #endif // BUILD_PARAMETERS_HPP diff --git a/include/constant_medium.hpp b/include/constant_medium.hpp deleted file mode 100644 index ebc8fc7..0000000 --- a/include/constant_medium.hpp +++ /dev/null @@ -1,86 +0,0 @@ -#ifndef CONSTANT_MEDIUM_HPP -#define CONSTANT_MEDIUM_HPP - -#include "box.hpp" -#include "material.hpp" -#include "sphere.hpp" -#include "texture.hpp" -#include "visit.hpp" - -using hittableVolume_t = std::variant; - -/** - * A ray going through the volume can either make it all the way through - * or be scattered at some point on or inside the volume. - */ -class constant_medium { - public: - constant_medium(const hittableVolume_t& b, real_t d, texture_t& a) - : boundary { b } - , neg_inv_density { -1 / d } - , phase_function { isotropic_material { a } } {} - - constant_medium(const hittableVolume_t& b, real_t d, const color& a) - : boundary { b } - , neg_inv_density { -1 / d } - , phase_function { isotropic_material { a } } {} - - bool hit(auto& ctx, const ray& r, real_t min, real_t max, hit_record& rec, - material_t& hit_material_type) const { - auto& rng = ctx.rng; - hit_material_type = phase_function; - material_t temp_material_type; - hit_record rec1, rec2; - if (!dev_visit(monostate_dispatch( - [&](auto&& object) { - return object.hit(ctx, r, -infinity, infinity, rec1, - temp_material_type); - }, - false), - boundary)) { - return false; - } - - if (!dev_visit(monostate_dispatch( - [&](auto&& object) { - return object.hit(ctx, r, rec1.t + 0.0001f, infinity, - rec2, temp_material_type); - }, - false), - boundary)) { - return false; - } - - if (rec1.t < min) - rec1.t = min; - if (rec2.t > max) - rec2.t = max; - if (rec1.t >= rec2.t) - return false; - if (rec1.t < 0) - rec1.t = 0; - - const auto ray_length = sycl::length(r.direction()); - /// Distance between the two hitpoints affect of probability - /// of the ray hitting a smoke particle - const auto distance_inside_boundary = (rec2.t - rec1.t) * ray_length; - const auto hit_distance = neg_inv_density * sycl::log(rng.real()); - - /// With lower density, hit_distance has higher probabilty - /// of being greater than distance_inside_boundary - if (hit_distance > distance_inside_boundary) - return false; - - rec.t = rec1.t + hit_distance / ray_length; - rec.p = r.at(rec.t); - - rec.normal = vec { 1, 0, 0 }; // arbitrary - rec.front_face = true; // also arbitrary - return true; - } - - hittableVolume_t boundary; - real_t neg_inv_density; - material_t phase_function; -}; -#endif diff --git a/include/hitable.hpp b/include/hitable.hpp index dfd8d39..4bfff59 100644 --- a/include/hitable.hpp +++ b/include/hitable.hpp @@ -1,13 +1,9 @@ #ifndef HITTABLE_H #define HITTABLE_H -#include "box.hpp" -#include "constant_medium.hpp" #include "ray.hpp" -#include "rectangle.hpp" #include "sphere.hpp" #include "triangle.hpp" -using hittable_t = std::variant; +using hittable_t = std::variant; #endif diff --git a/include/material.hpp b/include/material.hpp index a999fc6..13196fa 100644 --- a/include/material.hpp +++ b/include/material.hpp @@ -32,68 +32,6 @@ struct lambertian_material { texture_t albedo; }; -struct metal_material { - metal_material() = default; - metal_material(const color& a, real_t f) - : albedo { a } - , fuzz { std::clamp(f, 0.0f, 1.0f) } {} - - bool scatter(auto& ctx, const ray& r_in, const hit_record& rec, - color& attenuation, ray& scattered) const { - auto& rng = ctx.rng; - vec reflected = reflect(unit_vector(r_in.direction()), rec.normal); - scattered = ray(rec.p, reflected + fuzz * rng.in_unit_ball(), r_in.time()); - // Attenuation of the ray hitting the object is modified based on the color - // at hit point - attenuation *= albedo; - return (dot(scattered.direction(), rec.normal) > 0); - } - - color emitted(auto&, const hit_record& rec) { return color(0.f, 0.f, 0.f); } - color albedo; - real_t fuzz; -}; - -struct dielectric_material { - dielectric_material() = default; - dielectric_material(real_t ri, const color& albedo) - : ref_idx { ri } - , albedo { albedo } {} - - // Schlick's approximation for reflectance - real_t reflectance(real_t cosine, real_t ref_idx) const { - auto r0 = (1 - ref_idx) / (1 + ref_idx); - r0 *= r0; - return r0 + (1 - r0) * sycl::pow((1 - cosine), 5.0f); - } - - bool scatter(auto& ctx, const ray& r_in, const hit_record& rec, - color& attenuation, ray& scattered) const { - // Attenuation of the ray hitting the object is modified based on the color - // at hit point - auto& rng = ctx.rng; - attenuation *= albedo; - real_t refraction_ratio = rec.front_face ? (1.0f / ref_idx) : ref_idx; - vec unit_direction = unit_vector(r_in.direction()); - real_t cos_theta = sycl::fmin(-sycl::dot(unit_direction, rec.normal), 1.0f); - real_t sin_theta = sycl::sqrt(1.0f - cos_theta * cos_theta); - bool cannot_refract = refraction_ratio * sin_theta > 1.0f; - vec direction; - if (cannot_refract || reflectance(cos_theta, refraction_ratio) > rng.real()) - direction = reflect(unit_direction, rec.normal); - else - direction = refract(unit_direction, rec.normal, refraction_ratio); - - scattered = ray(rec.p, direction, r_in.time()); - return true; - } - - color emitted(auto&, const hit_record& rec) { return color(0, 0, 0); } - // Refractive index of the glass - real_t ref_idx; - // Color of the glass - color albedo; -}; struct lightsource_material { lightsource_material() = default; @@ -137,6 +75,5 @@ struct isotropic_material { }; using material_t = - std::variant; + std::variant; #endif diff --git a/include/rectangle.hpp b/include/rectangle.hpp deleted file mode 100644 index 7ac9a08..0000000 --- a/include/rectangle.hpp +++ /dev/null @@ -1,132 +0,0 @@ -#ifndef RECT_HPP -#define RECT_HPP - -#include "material.hpp" -#include "ray.hpp" -#include "rtweekend.hpp" -#include "texture.hpp" -#include "vec.hpp" - -/** The Following classes implement: - - - - https://raytracing.github.io/books/RayTracingTheNextWeek.html#rectanglesandlights/creatingrectangleobjectsa -*/ - -class xy_rect { - public: - xy_rect() = default; - - /// x0 <= x1 and y0 <= y1 - xy_rect(real_t _x0, real_t _x1, real_t _y0, real_t _y1, real_t _k, - const material_t& mat_type) - : x0 { _x0 } - , x1 { _x1 } - , y0 { _y0 } - , y1 { _y1 } - , k { _k } - , material_type { mat_type } {} - - /// Compute ray interaction with rectangle - bool hit(auto&, const ray& r, real_t min, real_t max, hit_record& rec, - material_t& hit_material_type) const { - hit_material_type = material_type; - - auto t = (k - r.origin().z()) / r.direction().z(); - if (t < min || t > max) - return false; - auto x = r.origin().x() + t * r.direction().x(); - auto y = r.origin().y() + t * r.direction().y(); - if (x < x0 || x > x1 || y < y0 || y > y1) - return false; - rec.u = (x - x0) / (x1 - x0); - rec.v = (y - y0) / (y1 - y0); - rec.t = t; - rec.p = r.at(rec.t); - vec outward_normal = vec(0, 0, 1); - rec.set_face_normal(r, outward_normal); - return true; - } - real_t x0, x1, y0, y1, k; - material_t material_type; -}; - -class xz_rect { - public: - xz_rect() = default; - - /// x0 <= x1 and z0 <= z1 - xz_rect(real_t _x0, real_t _x1, real_t _z0, real_t _z1, real_t _k, - const material_t& mat_type) - : x0 { _x0 } - , x1 { _x1 } - , z0 { _z0 } - , z1 { _z1 } - , k { _k } - , material_type { mat_type } {} - - /// Compute ray interaction with rectangle - bool hit(auto&, const ray& r, real_t min, real_t max, hit_record& rec, - material_t& hit_material_type) const { - hit_material_type = material_type; - - auto t = (k - r.origin().y()) / r.direction().y(); - if (t < min || t > max) - return false; - auto x = r.origin().x() + t * r.direction().x(); - auto z = r.origin().z() + t * r.direction().z(); - if (x < x0 || x > x1 || z < z0 || z > z1) - return false; - rec.u = (x - x0) / (x1 - x0); - rec.v = (z - z0) / (z1 - z0); - rec.t = t; - rec.p = r.at(rec.t); - vec outward_normal = vec(0, 1, 0); - rec.set_face_normal(r, outward_normal); - return true; - } - real_t x0, x1, z0, z1, k; - material_t material_type; -}; - -class yz_rect { - public: - yz_rect() = default; - - /// y0 <= y1 and z0 <= z1 - yz_rect(real_t _y0, real_t _y1, real_t _z0, real_t _z1, real_t _k, - const material_t& mat_type) - : y0 { _y0 } - , y1 { _y1 } - , z0 { _z0 } - , z1 { _z1 } - , k { _k } - , material_type { mat_type } {} - - /// Compute ray interaction with rectangle - bool hit(auto&, const ray& r, real_t min, real_t max, hit_record& rec, - material_t& hit_material_type) const { - hit_material_type = material_type; - - auto t = (k - r.origin().x()) / r.direction().x(); - if (t < min || t > max) - return false; - auto y = r.origin().y() + t * r.direction().y(); - auto z = r.origin().z() + t * r.direction().z(); - if (y < y0 || y > y1 || z < z0 || z > z1) - return false; - rec.u = (y - y0) / (y1 - y0); - rec.v = (z - z0) / (z1 - z0); - rec.t = t; - rec.p = r.at(rec.t); - vec outward_normal = vec(1, 0, 0); - rec.set_face_normal(r, outward_normal); - return true; - } - real_t y0, y1, z0, z1, k; - material_t material_type; -}; - -using rectangle_t = std::variant; - -#endif diff --git a/include/render.hpp b/include/render.hpp index 0713873..ecd77ab 100644 --- a/include/render.hpp +++ b/include/render.hpp @@ -106,11 +106,11 @@ struct PixelRender; template inline void executor(sycl::handler& cgh, camera const& cam_ptr, - auto& hittable_acc, auto& fb_acc, auto& texture_acc) { + auto& hittable_acc, auto& fb_acc) { if constexpr (buildparams::use_single_task) { cgh.single_task([=] { LocalPseudoRNG rng; - task_context ctx { rng, texture_acc.get_pointer() }; + task_context ctx { rng }; for (int x_coord = 0; x_coord != width; ++x_coord) for (int y_coord = 0; y_coord != height; ++y_coord) { render_pixel( @@ -127,7 +127,7 @@ inline void executor(sycl::handler& cgh, camera const& cam_ptr, auto init_generator_state = std::hash {}(item.get_linear_id()); LocalPseudoRNG rng(init_generator_state); - task_context ctx { rng, texture_acc.get_pointer() }; + task_context ctx { rng }; render_pixel( ctx, x_coord, y_coord, cam_ptr, hittable_acc, fb_acc); }); @@ -142,16 +142,12 @@ void render(sycl::queue& queue, sycl::buffer& frame_buf, const auto nb_hittable = hittables.size(); auto hittables_buf = sycl::buffer(hittables.data(), sycl::range<1>(nb_hittable)); - auto texture_buf = image_texture::freeze(); // Submit command group on device queue.submit([&](sycl::handler& cgh) { auto fb_acc = frame_buf.get_access(cgh); auto hittables_acc = hittables_buf.get_access(cgh); - auto texture_acc = texture_buf.get_access(cgh); - - executor(cgh, cam, hittables_acc, fb_acc, - texture_acc); + executor(cgh, cam, hittables_acc, fb_acc); }); } diff --git a/include/rtweekend.hpp b/include/rtweekend.hpp index cf0871d..9a0a248 100644 --- a/include/rtweekend.hpp +++ b/include/rtweekend.hpp @@ -102,8 +102,6 @@ class LocalPseudoRNG { */ struct task_context { LocalPseudoRNG rng; - // See image_texture in texture.hpp for more details - sycl::global_ptr texture_data; }; // Common Headers diff --git a/include/texture.hpp b/include/texture.hpp index 3753f98..d9abd0d 100644 --- a/include/texture.hpp +++ b/include/texture.hpp @@ -51,109 +51,6 @@ struct checker_texture { solid_texture even; }; -/** - @brief A texture based on an image +using texture_t = std::variant; - In order to be able to get the bitmap on the device without embedding it in - the object, all image_texture textures are serialized in one vector. - - The offset of the texture in the vector is stored in the image_texture - instance. - - When all the textures have been loaded, the freeze() method can be called to - get a sycl::buffer that store this data. - - */ -struct image_texture { - private: - static constexpr auto bytes_per_pixel = 3; - // Vector in which all the textures are serialized - static std::vector texture_data; - static bool frozen; - - std::size_t width {}; - std::size_t height {}; - // offset of the first pixel in the texture vector - std::size_t offset; - - /// The repetition rate of the image - real_t cyclic_frequency { 1.f }; - - image_texture(std::size_t _width, std::size_t _height, std::size_t _offset, - real_t _cyclic_frequency) - : width { _width } - , height { _height } - , offset { _offset } - , cyclic_frequency { _cyclic_frequency } {} - - public: - /** Create a texture from an image file - - \param[in] file_name is the path name to the image file - - \param[in] cyclic_frequency is an optional repetition rate of - the image in the texture - */ - static image_texture image_texture_factory(const char* file_name, - real_t _cyclic_frequency = 1) { - assert(!frozen); - auto components_per_pixel = bytes_per_pixel; - uint8_t* _data; - int _w, _h; - _data = - stbi_load(file_name, &_w, &_h, &components_per_pixel, bytes_per_pixel); - std::size_t _offset = 0; - if (!_data) { - std::cerr << "ERROR: Could not load texture image file '" << file_name - << "'.\n" - << stbi_failure_reason() << std::endl; - _w = _h = 1; - } else { - auto size = bytes_per_pixel * _w * _h; - _offset = texture_data.size() / 3; - std::copy(_data, _data + size, std::back_inserter(texture_data)); - } - return image_texture(_w, _h, _offset, _cyclic_frequency); - } - - /** - @brief Get a sycl::buffer containing texture data. - - image_texture_factory should not be called after having called freeze - - @return sycl::buffer - */ - static sycl::buffer freeze() { - assert(!frozen); - frozen = true; - return sycl::buffer { texture_data.data(), - { texture_data.size() / 3, 3 } }; - } - - /// Get the color for the texture at the given place - /// \todo rename this value() to color() everywhere? - color value(auto& ctx, const hit_record& rec) const { - // If texture data is unavailable, return solid cyan - // The image is repeated by the repetition factor - - std::size_t i = - sycl::fmod(rec.u * cyclic_frequency, (real_t)1) * (width - 1); - // The image frame buffer is going downwards, so flip the y axis - std::size_t j = - (1 - sycl::fmod(rec.v * cyclic_frequency, (real_t)1)) * (height - 1); - std::size_t local_offset = j * width + i; - std::size_t pix_idx = local_offset + offset; - auto scale = 1.f / 255; - auto& texture_data = ctx.texture_data; - return { texture_data[pix_idx * 3] * scale, - texture_data[pix_idx * 3 + 1] * scale, - texture_data[pix_idx * 3 + 2] * scale }; - } -}; - -using texture_t = std::variant; - -// Start filled with the fallback texture (solid blue) for texture load error -std::vector image_texture::texture_data { 0, 0, 1 }; -bool image_texture::frozen = false; #endif diff --git a/src/main.cpp b/src/main.cpp index 2456c25..36fee59 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -30,7 +30,7 @@ void dump_image_ppm(int width, int height, auto& fb_data) { } } -void save_image_png(int width, int height, sycl::buffer &fb) { +void save_image_png(int width, int height, sycl::buffer& fb) { constexpr unsigned num_channels = 3; auto fb_data = fb.get_access(); @@ -77,38 +77,18 @@ int main() { for (int a = -11; a < 11; a++) { for (int b = -11; b < 11; b++) { - // Based on a random variable , the material type is chosen - auto choose_mat = rng.real(); // Spheres are placed at a point randomly displaced from a,b point center(a + 0.9f * rng.real(), 0.2f, b + 0.9f * rng.real()); if (sycl::length((center - point(4, 0.2f, 0))) > 0.9f) { - if (choose_mat < 0.4f) { - // Lambertian - auto albedo = rng.vec_t() * rng.vec_t(); - hittables.emplace_back( - sphere(center, 0.2f, lambertian_material(albedo))); - } else if (choose_mat < 0.8f) { - // Lambertian movig spheres - auto albedo = rng.vec_t() * rng.vec_t(); - auto center2 = center + point { 0, rng.real(0, 0.25f), 0 }; - hittables.emplace_back(sphere(center, center2, 0.0f, 1.0f, 0.2f, - lambertian_material(albedo))); - } else if (choose_mat < 0.95f) { - // metal - auto albedo = rng.vec_t(0.5f, 1); - auto fuzz = rng.real(0, 0.5f); - hittables.emplace_back( - sphere(center, 0.2f, metal_material(albedo, fuzz))); - } else { - // glass - hittables.emplace_back( - sphere(center, 0.2f, - dielectric_material(1.5f, color { 1.0f, 1.0f, 1.0f }))); - } + // Lambertian + auto albedo = rng.vec_t() * rng.vec_t(); + hittables.emplace_back( + sphere(center, 0.2f, lambertian_material(albedo))); } } } + /* // Pyramid hittables.emplace_back( triangle(point { 6.5f, 0.0f, 1.30f }, point { 6.25f, 0.50f, 1.05f }, @@ -159,9 +139,9 @@ int main() { lambertian_material { color { 0.75f, 0.75f, 0.75f } } }; hittables.emplace_back( constant_medium { smoke_sphere, 1, color { 1, 1, 1 } }); - - // SYCL queue - sycl::queue myQueue; + */ + // SYCL queue + sycl::queue myQueue; // Camera setup /// Position of the camera @@ -183,7 +163,7 @@ int main() { }; // Sample per pixel - constexpr auto samples = 100; + constexpr auto samples = buildparams::samples; // SYCL render kernel From 8d3c2259f6e29a9909d477b1a4bb05a81b5e6f45 Mon Sep 17 00:00:00 2001 From: Luc Forget Date: Mon, 17 May 2021 08:19:15 -0700 Subject: [PATCH 05/14] Get rid of context, makeall dimension runtime values --- CMakeLists.txt | 21 --------------- include/build_parameters.hpp | 3 ++- include/camera.hpp | 2 +- include/material.hpp | 33 +++++++++++------------ include/render.hpp | 51 +++++++++++++++--------------------- include/rtweekend.hpp | 9 +++++++ include/sphere.hpp | 2 +- include/texture.hpp | 8 +++--- include/triangle.hpp | 2 +- src/main.cpp | 21 ++++++++------- 10 files changed, 67 insertions(+), 85 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 3d0cc4f..ffc030e 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -41,24 +41,6 @@ if(NOT CMAKE_BUILD_TYPE AND NOT CMAKE_CONFIGURATION_TYPES) "Debug" "Release" "MinSizeRel" "RelWithDebInfo") endif() -if(NOT OUTPUT_WIDTH) - message(STATUS "Setting output width to 800 as none was specified.") - set(OUTPUT_WIDTH "800" CACHE - STRING "Image width in pixel" FORCE) -endif() - -if(NOT OUTPUT_HEIGHT) - message(STATUS "Setting output height to 480 as none was specified.") - set(OUTPUT_HEIGHT "480" CACHE - STRING "Image height in pixel" FORCE) -endif() - -if(NOT SAMPLES) - message(STATUS "Setting samples to 32 as none was specified.") - set(SAMPLES "32" CACHE - STRING "Number of ray that are casted for each pixel" FORCE) -endif() - set(SYCL_RT_SRC_DIR ${CMAKE_CURRENT_SOURCE_DIR}/src) set(SYCL_RT_INCLUDE_DIR ${CMAKE_CURRENT_SOURCE_DIR}/include) @@ -66,9 +48,6 @@ set(SYCL_RT_INCLUDE_DIR ${CMAKE_CURRENT_SOURCE_DIR}/include) add_executable(sycl-rt ${SYCL_RT_SRC_DIR}/main.cpp) target_include_directories(sycl-rt PRIVATE ${SYCL_RT_INCLUDE_DIR}) -target_compile_definitions(sycl-rt PRIVATE OUTPUT_WIDTH=${OUTPUT_WIDTH}) -target_compile_definitions(sycl-rt PRIVATE OUTPUT_HEIGHT=${OUTPUT_HEIGHT}) -target_compile_definitions(sycl-rt PRIVATE SAMPLES=${SAMPLES}) # This is a SYCL program if ("${SYCL_CXX_COMPILER}" STREQUAL "") diff --git a/include/build_parameters.hpp b/include/build_parameters.hpp index 84d8b03..52054eb 100644 --- a/include/build_parameters.hpp +++ b/include/build_parameters.hpp @@ -14,9 +14,10 @@ constexpr bool use_sycl_compiler = USE_SYCL_COMPILER; constexpr bool use_sycl_compiler = false; #endif -constexpr int output_width = OUTPUT_WIDTH; +/*constexpr int output_width = OUTPUT_WIDTH; constexpr int output_height = OUTPUT_HEIGHT; constexpr int samples = SAMPLES; +constexpr int depth = DEPTH;//*/ } // namespace buildparams #endif // BUILD_PARAMETERS_HPP diff --git a/include/camera.hpp b/include/camera.hpp index 8265bcf..f67ddd5 100644 --- a/include/camera.hpp +++ b/include/camera.hpp @@ -96,7 +96,7 @@ class camera { return { origin + offset, lower_left_corner + s * horizontal + t * vertical - origin - offset, - rng.real(time0, time1) }; + time0 }; } }; diff --git a/include/material.hpp b/include/material.hpp index 13196fa..02b617a 100644 --- a/include/material.hpp +++ b/include/material.hpp @@ -15,24 +15,23 @@ struct lambertian_material { lambertian_material(const texture_t& a) : albedo { a } {} - bool scatter(auto& ctx, const ray& r_in, const hit_record& rec, + bool scatter(const ray& r_in, const hit_record& rec, color& attenuation, ray& scattered) const { - auto& rng = ctx.rng; + auto rng = LocalPseudoRNG { toseed(rec.normal) }; vec scatter_direction = rec.normal + rng.unit_vec(); scattered = ray(rec.p, scatter_direction, r_in.time()); // Attenuation of the ray hitting the object is modified based on the color // at hit point attenuation *= dev_visit( - monostate_dispatch([&](auto&& t) { return t.value(ctx, rec); }, - color { 0.f, 0.f, 0.f }), + monostate_dispatch([&](auto&& t) { return t.value(rec); }, + color { 0.f, 0.f, 0.f }), albedo); return true; } - color emitted(auto&, const hit_record& rec) { return color(0.f, 0.f, 0.f); } + color emitted(const hit_record& rec) { return color(0.f, 0.f, 0.f); } texture_t albedo; }; - struct lightsource_material { lightsource_material() = default; lightsource_material(const texture_t& a) @@ -42,9 +41,9 @@ struct lightsource_material { template bool scatter(T&...) const { return false; } - color emitted(auto& ctx, const hit_record& rec) { + color emitted(const hit_record& rec) { return dev_visit( - monostate_dispatch([&](auto&& t) { return t.value(ctx, rec); }, + monostate_dispatch([&](auto&& t) { return t.value(rec); }, color { 0.f, 0.f, 0.f }), emit); } @@ -58,22 +57,22 @@ struct isotropic_material { isotropic_material(texture_t& a) : albedo { a } {} - bool scatter(auto& ctx, const ray& r_in, const hit_record& rec, + bool scatter(const ray& r_in, const hit_record& rec, color& attenuation, ray& scattered) const { - auto& rng = ctx.rng; + LocalPseudoRNG rng(toseed(r_in.direction())); scattered = ray(rec.p, rng.in_unit_ball(), r_in.time()); - attenuation *= - dev_visit(monostate_dispatch([&](auto&& t) { return t.value(ctx, rec); }, - color { 0.f, 0.f, 0.f }), - albedo); + attenuation *= dev_visit( + monostate_dispatch([&](auto&& t) { return t.value(rec); }, + color { 0.f, 0.f, 0.f }), + albedo); return true; } - color emitted(auto&, const hit_record& rec) { return color(0, 0, 0); } + color emitted(const hit_record& rec) { return color(0, 0, 0); } texture_t albedo; }; -using material_t = - std::variant; +using material_t = std::variant; #endif diff --git a/include/render.hpp b/include/render.hpp index ecd77ab..32bac60 100644 --- a/include/render.hpp +++ b/include/render.hpp @@ -14,10 +14,8 @@ #include "vec.hpp" #include "visit.hpp" -template -inline auto render_pixel(auto& ctx, int x_coord, int y_coord, camera const& cam, - auto& hittable_acc, auto fb_acc) { - auto& rng = ctx.rng; +inline auto render_pixel(int width, int height, int depth, int samples, int x_coord, int y_coord, + camera const& cam, auto& hittable_acc, auto fb_acc) { auto get_color = [&](const ray& r) { auto hit_world = [&](const ray& r, hit_record& rec, material_t& material_type) { @@ -29,7 +27,7 @@ inline auto render_pixel(auto& ctx, int x_coord, int y_coord, camera const& cam, for (auto i = 0; i < hittable_acc.get_count(); i++) { if (dev_visit(monostate_dispatch( [&](auto&& object) { - return object.hit(ctx, r, 0.001f, closest_so_far, + return object.hit(r, 0.001f, closest_so_far, temp_rec, temp_material_type); }, false), @@ -51,15 +49,14 @@ inline auto render_pixel(auto& ctx, int x_coord, int y_coord, camera const& cam, for (auto i = 0; i < depth; i++) { hit_record rec; if (hit_world(cur_ray, rec, material_type)) { - emitted = - dev_visit(monostate_dispatch( - [&](auto&& mat) { return mat.emitted(ctx, rec); }, - color { 0.f, 0.f, 0.f }), - material_type); + emitted = dev_visit( + monostate_dispatch([&](auto&& mat) { return mat.emitted(rec); }, + color { 0.f, 0.f, 0.f }), + material_type); if (dev_visit(monostate_dispatch( [&](auto&& mat) { - return mat.scatter(ctx, cur_ray, rec, - cur_attenuation, scattered); + return mat.scatter(cur_ray, rec, cur_attenuation, + scattered); }, false), material_type)) { @@ -87,7 +84,8 @@ inline auto render_pixel(auto& ctx, int x_coord, int y_coord, camera const& cam, // If not returned within max_depth return black return color { 0.0f, 0.0f, 0.0f }; }; - + uint32_t seed = std::hash {}((x_coord << 16) | (y_coord & 0xFFFF)); + LocalPseudoRNG rng(seed); color final_color(0.0f, 0.0f, 0.0f); for (auto i = 0; i < samples; i++) { const auto u = (x_coord + rng.real()) / width; @@ -104,17 +102,14 @@ inline auto render_pixel(auto& ctx, int x_coord, int y_coord, camera const& cam, struct PixelRender; -template -inline void executor(sycl::handler& cgh, camera const& cam_ptr, - auto& hittable_acc, auto& fb_acc) { +inline void executor(int width, int height, int depth, int samples, sycl::handler& cgh, + camera const& cam_ptr, auto& hittable_acc, auto& fb_acc) { if constexpr (buildparams::use_single_task) { cgh.single_task([=] { - LocalPseudoRNG rng; - task_context ctx { rng }; for (int x_coord = 0; x_coord != width; ++x_coord) for (int y_coord = 0; y_coord != height; ++y_coord) { - render_pixel( - ctx, x_coord, y_coord, cam_ptr, hittable_acc, fb_acc); + render_pixel(width, height, depth, samples, x_coord, y_coord, + cam_ptr, hittable_acc, fb_acc); } }); } else { @@ -124,21 +119,16 @@ inline void executor(sycl::handler& cgh, camera const& cam_ptr, auto gid = item.get_id(); const auto x_coord = gid[1]; const auto y_coord = gid[0]; - auto init_generator_state = - std::hash {}(item.get_linear_id()); - LocalPseudoRNG rng(init_generator_state); - task_context ctx { rng }; - render_pixel( - ctx, x_coord, y_coord, cam_ptr, hittable_acc, fb_acc); + render_pixel(width, height, depth, samples, x_coord, y_coord, + cam_ptr, hittable_acc, fb_acc); }); } } // Render function to call the render kernel -template -void render(sycl::queue& queue, sycl::buffer& frame_buf, +void render(int width, int height, int depth, int samples, sycl::queue& queue, + sycl::buffer& frame_buf, std::vector& hittables, camera& cam) { - auto constexpr depth = 50; const auto nb_hittable = hittables.size(); auto hittables_buf = sycl::buffer(hittables.data(), sycl::range<1>(nb_hittable)); @@ -148,6 +138,7 @@ void render(sycl::queue& queue, sycl::buffer& frame_buf, auto fb_acc = frame_buf.get_access(cgh); auto hittables_acc = hittables_buf.get_access(cgh); - executor(cgh, cam, hittables_acc, fb_acc); + executor(width, height, depth, samples, cgh, cam, hittables_acc, + fb_acc); }); } diff --git a/include/rtweekend.hpp b/include/rtweekend.hpp index 9a0a248..fb14e6a 100644 --- a/include/rtweekend.hpp +++ b/include/rtweekend.hpp @@ -4,6 +4,7 @@ #include #include #include +#include #include #include #include @@ -34,6 +35,14 @@ using vec = real_vec; inline real_t degrees_to_radians(real_t degrees) { return degrees * pi / 180.0f; } +uint32_t toseed(vec const & val) { + uint32_t x, y, z; + std::memcpy(&x, &val.x(), sizeof(uint32_t)); + std::memcpy(&y, &val.y(), sizeof(uint32_t)); + std::memcpy(&z, &val.z(), sizeof(uint32_t)); + return x*x*y*y*z*z; +} + class LocalPseudoRNG { public: inline LocalPseudoRNG(std::uint32_t init_state = xorshift<>::initial_state) diff --git a/include/sphere.hpp b/include/sphere.hpp index 33624f1..804bb02 100644 --- a/include/sphere.hpp +++ b/include/sphere.hpp @@ -56,7 +56,7 @@ class sphere { } /// Compute ray interaction with sphere - bool hit(auto&, const ray& r, real_t min, real_t max, hit_record& rec, + bool hit(const ray& r, real_t min, real_t max, hit_record& rec, material_t& hit_material_type) const { hit_material_type = material_type; diff --git a/include/texture.hpp b/include/texture.hpp index d9abd0d..46786c6 100644 --- a/include/texture.hpp +++ b/include/texture.hpp @@ -22,7 +22,7 @@ struct solid_texture { solid_texture(real_t red, real_t green, real_t blue) : solid_texture { color { red, green, blue } } {} // For solid texture, the color is same throughout the sphere - color value(auto&, const hit_record&) const { return color_value; } + color value(const hit_record&) const { return color_value; } private: color color_value; @@ -39,13 +39,13 @@ struct checker_texture { : odd { solid_texture { c1 } } , even { solid_texture { c2 } } {} // Color value is different based on normalised spherical coordinates - color value(auto& ctx, const hit_record& rec) const { + color value(const hit_record& rec) const { auto sines = sycl::sin(10 * rec.p.x()) * sycl::sin(10 * rec.p.y()) * sycl::sin(10 * rec.p.z()); if (sines < 0) - return odd.value(ctx, rec); + return odd.value(rec); else - return even.value(ctx, rec); + return even.value(rec); } solid_texture odd; solid_texture even; diff --git a/include/triangle.hpp b/include/triangle.hpp index b1102c8..c8c0668 100644 --- a/include/triangle.hpp +++ b/include/triangle.hpp @@ -110,7 +110,7 @@ class _triangle : public _triangle_coord { , material_type { mat_type } {} /// Compute ray interaction with triangle - bool hit(auto&, const ray& r, real_t min, real_t max, hit_record& rec, + bool hit(const ray& r, real_t min, real_t max, hit_record& rec, material_t& hit_material_type) const { hit_material_type = material_type; return IntersectionStrategy(r, *this, min, max, rec); diff --git a/src/main.cpp b/src/main.cpp index 36fee59..4d52425 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -58,10 +58,16 @@ void save_image_png(int width, int height, sycl::buffer& fb) { width * num_channels); } -int main() { +int main(int argc, char* argv[]) { + if (argc != 5) { + std::cerr << "Usage: sycl-rt OUT_WIDTH OUT_HEIGHT DEPTH SAMPLES" << std::endl; + return -1; + } // Frame buffer dimensions - constexpr auto width = buildparams::output_width; - constexpr auto height = buildparams::output_height; + auto width = std::stoi({ argv[1] }); + auto height = std::stoi({ argv[2] }); + auto depth = std::stoi({ argv[3] }); + auto samples = std::stoi({ argv[4] }); /// Graphical objects std::vector hittables; @@ -75,8 +81,8 @@ int main() { LocalPseudoRNG rng; - for (int a = -11; a < 11; a++) { - for (int b = -11; b < 11; b++) { + for (int a = -11; a < 11; a+=3) { + for (int b = -11; b < 11; b+=3) { // Spheres are placed at a point randomly displaced from a,b point center(a + 0.9f * rng.real(), 0.2f, b + 0.9f * rng.real()); if (sycl::length((center - point(4, 0.2f, 0))) > 0.9f) { @@ -162,13 +168,10 @@ int main() { aperture, focus_dist, 0.0f, 1.0f }; - // Sample per pixel - constexpr auto samples = buildparams::samples; - // SYCL render kernel sycl::buffer fb(sycl::range<2>(height, width)); - render(myQueue, fb, hittables, cam); + render(width, height, depth, samples, myQueue, fb, hittables, cam); // Save image to file save_image_png(width, height, fb); From 6ceb4cecc3852bd1e8efdbb88468c19d5b27dbb0 Mon Sep 17 00:00:00 2001 From: Luc Forget Date: Tue, 25 May 2021 06:59:35 -0700 Subject: [PATCH 06/14] WIP feature restauration --- include/box.hpp | 58 ++++++++++++++++ include/constant_medium.hpp | 84 +++++++++++++++++++++++ include/rectangle.hpp | 132 ++++++++++++++++++++++++++++++++++++ 3 files changed, 274 insertions(+) create mode 100644 include/box.hpp create mode 100644 include/constant_medium.hpp create mode 100644 include/rectangle.hpp diff --git a/include/box.hpp b/include/box.hpp new file mode 100644 index 0000000..e0bcb82 --- /dev/null +++ b/include/box.hpp @@ -0,0 +1,58 @@ +#ifndef BOX_HPP +#define BOX_HPP + +#include "rectangle.hpp" +#include "rtweekend.hpp" +#include "visit.hpp" + +/// This class implements a axis aligned cuboid using 6 rectangles +class box { + public: + box() = default; + + /// p0 = { x0, y0, z0 } and p1 = { x1, y1. z1 } + /// where x0 <= x1, y0 <= y1 and z0 <= z1 + box(const point& p0, const point& p1, const material_t& mat_type) + : box_min { p0 } + , box_max { p1 } + , material_type { mat_type } { + /// Add six sides of the box based on box_min and box_max to sides + sides[0] = xy_rect(p0.x(), p1.x(), p0.y(), p1.y(), p1.z(), mat_type); + sides[1] = xy_rect(p0.x(), p1.x(), p0.y(), p1.y(), p0.z(), mat_type); + sides[2] = xz_rect(p0.x(), p1.x(), p0.z(), p1.z(), p1.y(), mat_type); + sides[3] = xz_rect(p0.x(), p1.x(), p0.z(), p1.z(), p0.y(), mat_type); + sides[4] = yz_rect(p0.y(), p1.y(), p0.z(), p1.z(), p1.x(), mat_type); + sides[5] = yz_rect(p0.y(), p1.y(), p0.z(), p1.z(), p0.x(), mat_type); + } + + /// Compute ray interaction with the box + bool hit(const ray& r, real_t min, real_t max, hit_record& rec, + material_t& hit_material_type) const { + hit_record temp_rec; + material_t temp_material_type; + auto hit_anything = false; + auto closest_so_far = max; + // Checking if the ray hits any of the sides + for (const auto& side : sides) { + if (dev_visit( + [&](auto&& arg) { + return arg.hit(r, min, closest_so_far, temp_rec, + temp_material_type); + }, + side)) { + hit_anything = true; + closest_so_far = temp_rec.t; + rec = temp_rec; + hit_material_type = temp_material_type; + } + } + return hit_anything; + } + + point box_min; + point box_max; + material_t material_type; + std::array sides; +}; + +#endif diff --git a/include/constant_medium.hpp b/include/constant_medium.hpp new file mode 100644 index 0000000..26a02d1 --- /dev/null +++ b/include/constant_medium.hpp @@ -0,0 +1,84 @@ +#ifndef CONSTANT_MEDIUM_HPP +#define CONSTANT_MEDIUM_HPP + +#include "box.hpp" +#include "material.hpp" +#include "sphere.hpp" +#include "texture.hpp" +#include "visit.hpp" + +using hittableVolume_t = std::variant; + +/** + * A ray going through the volume can either make it all the way through + * or be scattered at some point on or inside the volume. + */ +class constant_medium { + public: + constant_medium(const hittableVolume_t& b, real_t d, texture_t& a) + : boundary { b } + , neg_inv_density { -1 / d } + , phase_function { isotropic_material { a } } {} + + constant_medium(const hittableVolume_t& b, real_t d, const color& a) + : boundary { b } + , neg_inv_density { -1 / d } + , phase_function { isotropic_material { a } } {} + + bool hit(const ray& r, real_t min, real_t max, hit_record& rec, + material_t& hit_material_type) const { + hit_material_type = phase_function; + material_t temp_material_type; + hit_record rec1, rec2; + if (!dev_visit( + [&](auto&& arg) { + return arg.hit(r, -infinity, infinity, rec1, + temp_material_type); + }, + boundary)) { + return false; + } + + if (!dev_visit( + [&](auto&& arg) { + return arg.hit(r, rec1.t + 0.0001f, infinity, rec2, + temp_material_type); + }, + boundary)) { + return false; + } + + if (rec1.t < min) + rec1.t = min; + if (rec2.t > max) + rec2.t = max; + if (rec1.t >= rec2.t) + return false; + if (rec1.t < 0) + rec1.t = 0; + + const auto ray_length = sycl::length(r.direction()); + /// Distance between the two hitpoints affect of probability + /// of the ray hitting a smoke particle + const auto distance_inside_boundary = (rec2.t - rec1.t) * ray_length; + auto rng = LocalPseudoRNG { toseed(r.direction()) }; + const auto hit_distance = neg_inv_density * sycl::log(rng.real()); + + /// With lower density, hit_distance has higher probabilty + /// of being greater than distance_inside_boundary + if (hit_distance > distance_inside_boundary) + return false; + + rec.t = rec1.t + hit_distance / ray_length; + rec.p = r.at(rec.t); + + rec.normal = vec { 1, 0, 0 }; // arbitrary + rec.front_face = true; // also arbitrary + return true; + } + + hittableVolume_t boundary; + real_t neg_inv_density; + material_t phase_function; +}; +#endif diff --git a/include/rectangle.hpp b/include/rectangle.hpp new file mode 100644 index 0000000..f36cbd2 --- /dev/null +++ b/include/rectangle.hpp @@ -0,0 +1,132 @@ +#ifndef RECT_HPP +#define RECT_HPP + +#include "material.hpp" +#include "ray.hpp" +#include "rtweekend.hpp" +#include "texture.hpp" +#include "vec.hpp" + +/** The Following classes implement: + + - + https://raytracing.github.io/books/RayTracingTheNextWeek.html#rectanglesandlights/creatingrectangleobjectsa +*/ + +class xy_rect { + public: + xy_rect() = default; + + /// x0 <= x1 and y0 <= y1 + xy_rect(real_t _x0, real_t _x1, real_t _y0, real_t _y1, real_t _k, + const material_t& mat_type) + : x0 { _x0 } + , x1 { _x1 } + , y0 { _y0 } + , y1 { _y1 } + , k { _k } + , material_type { mat_type } {} + + /// Compute ray interaction with rectangle + bool hit(const ray& r, real_t min, real_t max, hit_record& rec, + material_t& hit_material_type) const { + hit_material_type = material_type; + + auto t = (k - r.origin().z()) / r.direction().z(); + if (t < min || t > max) + return false; + auto x = r.origin().x() + t * r.direction().x(); + auto y = r.origin().y() + t * r.direction().y(); + if (x < x0 || x > x1 || y < y0 || y > y1) + return false; + rec.u = (x - x0) / (x1 - x0); + rec.v = (y - y0) / (y1 - y0); + rec.t = t; + rec.p = r.at(rec.t); + vec outward_normal = vec(0, 0, 1); + rec.set_face_normal(r, outward_normal); + return true; + } + real_t x0, x1, y0, y1, k; + material_t material_type; +}; + +class xz_rect { + public: + xz_rect() = default; + + /// x0 <= x1 and z0 <= z1 + xz_rect(real_t _x0, real_t _x1, real_t _z0, real_t _z1, real_t _k, + const material_t& mat_type) + : x0 { _x0 } + , x1 { _x1 } + , z0 { _z0 } + , z1 { _z1 } + , k { _k } + , material_type { mat_type } {} + + /// Compute ray interaction with rectangle + bool hit(const ray& r, real_t min, real_t max, hit_record& rec, + material_t& hit_material_type) const { + hit_material_type = material_type; + + auto t = (k - r.origin().y()) / r.direction().y(); + if (t < min || t > max) + return false; + auto x = r.origin().x() + t * r.direction().x(); + auto z = r.origin().z() + t * r.direction().z(); + if (x < x0 || x > x1 || z < z0 || z > z1) + return false; + rec.u = (x - x0) / (x1 - x0); + rec.v = (z - z0) / (z1 - z0); + rec.t = t; + rec.p = r.at(rec.t); + vec outward_normal = vec(0, 1, 0); + rec.set_face_normal(r, outward_normal); + return true; + } + real_t x0, x1, z0, z1, k; + material_t material_type; +}; + +class yz_rect { + public: + yz_rect() = default; + + /// y0 <= y1 and z0 <= z1 + yz_rect(real_t _y0, real_t _y1, real_t _z0, real_t _z1, real_t _k, + const material_t& mat_type) + : y0 { _y0 } + , y1 { _y1 } + , z0 { _z0 } + , z1 { _z1 } + , k { _k } + , material_type { mat_type } {} + + /// Compute ray interaction with rectangle + bool hit(const ray& r, real_t min, real_t max, hit_record& rec, + material_t& hit_material_type) const { + hit_material_type = material_type; + + auto t = (k - r.origin().x()) / r.direction().x(); + if (t < min || t > max) + return false; + auto y = r.origin().y() + t * r.direction().y(); + auto z = r.origin().z() + t * r.direction().z(); + if (y < y0 || y > y1 || z < z0 || z > z1) + return false; + rec.u = (y - y0) / (y1 - y0); + rec.v = (z - z0) / (z1 - z0); + rec.t = t; + rec.p = r.at(rec.t); + vec outward_normal = vec(1, 0, 0); + rec.set_face_normal(r, outward_normal); + return true; + } + real_t y0, y1, z0, z1, k; + material_t material_type; +}; + +using rectangle_t = std::variant; + +#endif From 60c883ab625172904188100b0e8d6f8cd67e3cc3 Mon Sep 17 00:00:00 2001 From: Luc Forget Date: Tue, 25 May 2021 07:01:08 -0700 Subject: [PATCH 07/14] Restore Features --- include/hitable.hpp | 4 +- include/material.hpp | 99 ++++++++++++++++++++++++++++++++++++-------- src/main.cpp | 21 +++++----- 3 files changed, 94 insertions(+), 30 deletions(-) diff --git a/include/hitable.hpp b/include/hitable.hpp index 4bfff59..4302399 100644 --- a/include/hitable.hpp +++ b/include/hitable.hpp @@ -4,6 +4,8 @@ #include "ray.hpp" #include "sphere.hpp" #include "triangle.hpp" +#include "box.hpp" +#include "constant_medium.hpp" -using hittable_t = std::variant; +using hittable_t = std::variant; #endif diff --git a/include/material.hpp b/include/material.hpp index 02b617a..b07364b 100644 --- a/include/material.hpp +++ b/include/material.hpp @@ -15,23 +15,86 @@ struct lambertian_material { lambertian_material(const texture_t& a) : albedo { a } {} - bool scatter(const ray& r_in, const hit_record& rec, - color& attenuation, ray& scattered) const { + bool scatter(const ray& r_in, const hit_record& rec, color& attenuation, + ray& scattered) const { auto rng = LocalPseudoRNG { toseed(rec.normal) }; vec scatter_direction = rec.normal + rng.unit_vec(); scattered = ray(rec.p, scatter_direction, r_in.time()); // Attenuation of the ray hitting the object is modified based on the color // at hit point - attenuation *= dev_visit( - monostate_dispatch([&](auto&& t) { return t.value(rec); }, - color { 0.f, 0.f, 0.f }), - albedo); + attenuation *= + dev_visit(monostate_dispatch([&](auto&& t) { return t.value(rec); }, + color { 0.f, 0.f, 0.f }), + albedo); return true; } color emitted(const hit_record& rec) { return color(0.f, 0.f, 0.f); } texture_t albedo; }; +struct metal_material { + metal_material() = default; + metal_material(const color& a, real_t f) + : albedo { a } + , fuzz { std::clamp(f, 0.0f, 1.0f) } {} + + bool scatter(const ray& r_in, const hit_record& rec, + color& attenuation, ray& scattered) const { + auto rng = LocalPseudoRNG { toseed(rec.normal) }; + vec reflected = reflect(unit_vector(r_in.direction()), rec.normal); + scattered = ray(rec.p, reflected + fuzz * rng.in_unit_ball(), r_in.time()); + // Attenuation of the ray hitting the object is modified based on the color + // at hit point + attenuation *= albedo; + return (dot(scattered.direction(), rec.normal) > 0); + } + + color emitted(const hit_record& rec) { return color(0, 0, 0); } + color albedo; + real_t fuzz; +}; + +struct dielectric_material { + dielectric_material() = default; + dielectric_material(real_t ri, const color& albedo) + : ref_idx { ri } + , albedo { albedo } {} + + // Schlick's approximation for reflectance + real_t reflectance(real_t cosine, real_t ref_idx) const { + auto r0 = (1 - ref_idx) / (1 + ref_idx); + r0 *= r0; + return r0 + (1 - r0) * sycl::pow((1 - cosine), 5.0f); + } + + bool scatter(const ray& r_in, const hit_record& rec, + color& attenuation, ray& scattered) const { + // Attenuation of the ray hitting the object is modified based on the color + // at hit point + auto rng = LocalPseudoRNG { toseed(rec.normal) }; + attenuation *= albedo; + real_t refraction_ratio = rec.front_face ? (1.0f / ref_idx) : ref_idx; + vec unit_direction = unit_vector(r_in.direction()); + real_t cos_theta = sycl::fmin(-sycl::dot(unit_direction, rec.normal), 1.0f); + real_t sin_theta = sycl::sqrt(1.0f - cos_theta * cos_theta); + bool cannot_refract = refraction_ratio * sin_theta > 1.0f; + vec direction; + if (cannot_refract || reflectance(cos_theta, refraction_ratio) > rng.real()) + direction = reflect(unit_direction, rec.normal); + else + direction = refract(unit_direction, rec.normal, refraction_ratio); + + scattered = ray(rec.p, direction, r_in.time()); + return true; + } + + color emitted(const hit_record& rec) { return color(0, 0, 0); } + // Refractive index of the glass + real_t ref_idx; + // Color of the glass + color albedo; +}; + struct lightsource_material { lightsource_material() = default; lightsource_material(const texture_t& a) @@ -42,10 +105,9 @@ struct lightsource_material { template bool scatter(T&...) const { return false; } color emitted(const hit_record& rec) { - return dev_visit( - monostate_dispatch([&](auto&& t) { return t.value(rec); }, - color { 0.f, 0.f, 0.f }), - emit); + return dev_visit(monostate_dispatch([&](auto&& t) { return t.value(rec); }, + color { 0.f, 0.f, 0.f }), + emit); } texture_t emit; @@ -57,14 +119,14 @@ struct isotropic_material { isotropic_material(texture_t& a) : albedo { a } {} - bool scatter(const ray& r_in, const hit_record& rec, - color& attenuation, ray& scattered) const { + bool scatter(const ray& r_in, const hit_record& rec, color& attenuation, + ray& scattered) const { LocalPseudoRNG rng(toseed(r_in.direction())); scattered = ray(rec.p, rng.in_unit_ball(), r_in.time()); - attenuation *= dev_visit( - monostate_dispatch([&](auto&& t) { return t.value(rec); }, - color { 0.f, 0.f, 0.f }), - albedo); + attenuation *= + dev_visit(monostate_dispatch([&](auto&& t) { return t.value(rec); }, + color { 0.f, 0.f, 0.f }), + albedo); return true; } @@ -73,6 +135,7 @@ struct isotropic_material { texture_t albedo; }; -using material_t = std::variant; +using material_t = + std::variant; #endif diff --git a/src/main.cpp b/src/main.cpp index 4d52425..9258bd7 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -40,7 +40,6 @@ void save_image_png(int width, int height, sycl::buffer& fb) { int index = 0; for (int j = height - 1; j >= 0; --j) { for (int i = 0; i < width; ++i) { - auto input_index = j * width + i; int r = static_cast( 256 * std::clamp(sycl::sqrt(fb_data[j][i].x()), 0.0f, 0.999f)); int g = static_cast( @@ -60,7 +59,8 @@ void save_image_png(int width, int height, sycl::buffer& fb) { int main(int argc, char* argv[]) { if (argc != 5) { - std::cerr << "Usage: sycl-rt OUT_WIDTH OUT_HEIGHT DEPTH SAMPLES" << std::endl; + std::cerr << "Usage: sycl-rt OUT_WIDTH OUT_HEIGHT DEPTH SAMPLES" + << std::endl; return -1; } // Frame buffer dimensions @@ -81,8 +81,8 @@ int main(int argc, char* argv[]) { LocalPseudoRNG rng; - for (int a = -11; a < 11; a+=3) { - for (int b = -11; b < 11; b+=3) { + for (int a = -11; a < 11; a += 3) { + for (int b = -11; b < 11; b += 3) { // Spheres are placed at a point randomly displaced from a,b point center(a + 0.9f * rng.real(), 0.2f, b + 0.9f * rng.real()); if (sycl::length((center - point(4, 0.2f, 0))) > 0.9f) { @@ -94,7 +94,6 @@ int main(int argc, char* argv[]) { } } - /* // Pyramid hittables.emplace_back( triangle(point { 6.5f, 0.0f, 1.30f }, point { 6.25f, 0.50f, 1.05f }, @@ -116,8 +115,8 @@ int main(int argc, char* argv[]) { sphere(point { 4, 1, 0 }, 0.2f, lightsource_material(color(10, 0, 10)))); // Four large spheres of metal, dielectric and Lambertian material types - t = image_texture::image_texture_factory("../images/Xilinx.jpg"); - hittables.emplace_back(xy_rect(2, 4, 0, 1, -1, lambertian_material(t))); + // t = image_texture::image_texture_factory("../images/Xilinx.jpg"); + //hittables.emplace_back(xy_rect(2, 4, 0, 1, -1, lambertian_material(t))); hittables.emplace_back( sphere(point { 4, 1, 2.25f }, 1, lambertian_material(t))); hittables.emplace_back( @@ -128,7 +127,7 @@ int main(int argc, char* argv[]) { hittables.emplace_back(sphere(point { 0, 1, -2.25f }, 1, metal_material(color(0.7f, 0.6f, 0.5f), 0.0f))); - t = image_texture::image_texture_factory("../images/SYCL.png", 5); + // t = image_texture::image_texture_factory("../images/SYCL.png", 5); // // Add a sphere with a SYCL logo in the background hittables.emplace_back( @@ -145,9 +144,9 @@ int main(int argc, char* argv[]) { lambertian_material { color { 0.75f, 0.75f, 0.75f } } }; hittables.emplace_back( constant_medium { smoke_sphere, 1, color { 1, 1, 1 } }); - */ - // SYCL queue - sycl::queue myQueue; + + // SYCL queue + sycl::queue myQueue; // Camera setup /// Position of the camera From 8a7354741654adc472f7283013b2bdaf4f74067d Mon Sep 17 00:00:00 2001 From: Luc Forget Date: Wed, 9 Jun 2021 11:53:20 -0700 Subject: [PATCH 08/14] Get rid of monostate --- include/constant_medium.hpp | 1 - include/hitable.hpp | 2 +- include/material.hpp | 14 ++++---------- include/render.hpp | 31 +++++++++++++------------------ include/texture.hpp | 2 +- src/main.cpp | 6 +++--- 6 files changed, 22 insertions(+), 34 deletions(-) diff --git a/include/constant_medium.hpp b/include/constant_medium.hpp index 26a02d1..236a988 100644 --- a/include/constant_medium.hpp +++ b/include/constant_medium.hpp @@ -38,7 +38,6 @@ class constant_medium { boundary)) { return false; } - if (!dev_visit( [&](auto&& arg) { return arg.hit(r, rec1.t + 0.0001f, infinity, rec2, diff --git a/include/hitable.hpp b/include/hitable.hpp index 4302399..63324d6 100644 --- a/include/hitable.hpp +++ b/include/hitable.hpp @@ -7,5 +7,5 @@ #include "box.hpp" #include "constant_medium.hpp" -using hittable_t = std::variant; +using hittable_t = std::variant; #endif diff --git a/include/material.hpp b/include/material.hpp index b07364b..e21dcb2 100644 --- a/include/material.hpp +++ b/include/material.hpp @@ -23,8 +23,7 @@ struct lambertian_material { // Attenuation of the ray hitting the object is modified based on the color // at hit point attenuation *= - dev_visit(monostate_dispatch([&](auto&& t) { return t.value(rec); }, - color { 0.f, 0.f, 0.f }), + dev_visit([&](auto&& t) { return t.value(rec); }, albedo); return true; } @@ -105,9 +104,7 @@ struct lightsource_material { template bool scatter(T&...) const { return false; } color emitted(const hit_record& rec) { - return dev_visit(monostate_dispatch([&](auto&& t) { return t.value(rec); }, - color { 0.f, 0.f, 0.f }), - emit); + return dev_visit([&](auto&& t) { return t.value(rec); }, emit); } texture_t emit; @@ -123,10 +120,7 @@ struct isotropic_material { ray& scattered) const { LocalPseudoRNG rng(toseed(r_in.direction())); scattered = ray(rec.p, rng.in_unit_ball(), r_in.time()); - attenuation *= - dev_visit(monostate_dispatch([&](auto&& t) { return t.value(rec); }, - color { 0.f, 0.f, 0.f }), - albedo); + attenuation *= dev_visit([&](auto&& t) { return t.value(rec); }, albedo); return true; } @@ -136,6 +130,6 @@ struct isotropic_material { }; using material_t = - std::variant; #endif diff --git a/include/render.hpp b/include/render.hpp index 32bac60..cf2787c 100644 --- a/include/render.hpp +++ b/include/render.hpp @@ -25,13 +25,12 @@ inline auto render_pixel(int width, int height, int depth, int samples, int x_co auto closest_so_far = infinity; // Checking if the ray hits any of the spheres for (auto i = 0; i < hittable_acc.get_count(); i++) { - if (dev_visit(monostate_dispatch( - [&](auto&& object) { - return object.hit(r, 0.001f, closest_so_far, - temp_rec, temp_material_type); - }, - false), - hittable_acc[i])) { + if (dev_visit( + [&](auto&& object) { + return object.hit(r, 0.001f, closest_so_far, temp_rec, + temp_material_type); + }, + hittable_acc[i])) { hit_anything = true; closest_so_far = temp_rec.t; rec = temp_rec; @@ -49,17 +48,13 @@ inline auto render_pixel(int width, int height, int depth, int samples, int x_co for (auto i = 0; i < depth; i++) { hit_record rec; if (hit_world(cur_ray, rec, material_type)) { - emitted = dev_visit( - monostate_dispatch([&](auto&& mat) { return mat.emitted(rec); }, - color { 0.f, 0.f, 0.f }), - material_type); - if (dev_visit(monostate_dispatch( - [&](auto&& mat) { - return mat.scatter(cur_ray, rec, cur_attenuation, - scattered); - }, - false), - material_type)) { + emitted = dev_visit([&](auto&& mat) { return mat.emitted(rec); }, + material_type); + if (dev_visit( + [&](auto&& mat) { + return mat.scatter(cur_ray, rec, cur_attenuation, scattered); + }, + material_type)) { // On hitting the object, the ray gets scattered cur_ray = scattered; } else { diff --git a/include/texture.hpp b/include/texture.hpp index 46786c6..cece438 100644 --- a/include/texture.hpp +++ b/include/texture.hpp @@ -51,6 +51,6 @@ struct checker_texture { solid_texture even; }; -using texture_t = std::variant; +using texture_t = std::variant; #endif diff --git a/src/main.cpp b/src/main.cpp index 9258bd7..dea2ea4 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -136,14 +136,14 @@ int main(int argc, char* argv[]) { // Add a metallic monolith hittables.emplace_back( box { point { 6.5f, 0, -1.5f }, point { 7.0f, 3.0f, -1.0f }, - metal_material { color { 0.7f, 0.6f, 0.5f }, 0.25f } }); + metal_material { color { 0.7f, 0.6f, 0.5f }, 0.25f } });/**/ // Add a smoke ball sphere smoke_sphere = sphere { point { 5, 1, 3.5f }, 1, lambertian_material { color { 0.75f, 0.75f, 0.75f } } }; - hittables.emplace_back( - constant_medium { smoke_sphere, 1, color { 1, 1, 1 } }); + /*hittables.emplace_back( + constant_medium { smoke_sphere, 1, color { 1, 1, 1 } });*/ // SYCL queue sycl::queue myQueue; From 4644777428f5ecdc7d8433a60baaeb5be40c632a Mon Sep 17 00:00:00 2001 From: Luc Forget Date: Fri, 11 Jun 2021 00:48:35 -0700 Subject: [PATCH 09/14] Formatting --- include/constant_medium.hpp | 3 +-- include/hit_record.hpp | 2 +- include/hitable.hpp | 4 ++-- include/material.hpp | 16 +++++++--------- include/render.hpp | 21 +++++++++++---------- include/rtweekend.hpp | 10 ++++++---- include/visit.hpp | 3 ++- 7 files changed, 30 insertions(+), 29 deletions(-) diff --git a/include/constant_medium.hpp b/include/constant_medium.hpp index 236a988..8ee509b 100644 --- a/include/constant_medium.hpp +++ b/include/constant_medium.hpp @@ -32,8 +32,7 @@ class constant_medium { hit_record rec1, rec2; if (!dev_visit( [&](auto&& arg) { - return arg.hit(r, -infinity, infinity, rec1, - temp_material_type); + return arg.hit(r, -infinity, infinity, rec1, temp_material_type); }, boundary)) { return false; diff --git a/include/hit_record.hpp b/include/hit_record.hpp index 342fb81..7b870dc 100644 --- a/include/hit_record.hpp +++ b/include/hit_record.hpp @@ -5,7 +5,7 @@ class hit_record { public: - real_t t; // + real_t t; // point p; // hit point vec normal; // normal at hit point bool front_face; // to check if hit point is on the outer surface diff --git a/include/hitable.hpp b/include/hitable.hpp index 63324d6..41a8ff1 100644 --- a/include/hitable.hpp +++ b/include/hitable.hpp @@ -1,11 +1,11 @@ #ifndef HITTABLE_H #define HITTABLE_H +#include "box.hpp" +#include "constant_medium.hpp" #include "ray.hpp" #include "sphere.hpp" #include "triangle.hpp" -#include "box.hpp" -#include "constant_medium.hpp" using hittable_t = std::variant; #endif diff --git a/include/material.hpp b/include/material.hpp index e21dcb2..0dfc648 100644 --- a/include/material.hpp +++ b/include/material.hpp @@ -22,9 +22,7 @@ struct lambertian_material { scattered = ray(rec.p, scatter_direction, r_in.time()); // Attenuation of the ray hitting the object is modified based on the color // at hit point - attenuation *= - dev_visit([&](auto&& t) { return t.value(rec); }, - albedo); + attenuation *= dev_visit([&](auto&& t) { return t.value(rec); }, albedo); return true; } color emitted(const hit_record& rec) { return color(0.f, 0.f, 0.f); } @@ -37,8 +35,8 @@ struct metal_material { : albedo { a } , fuzz { std::clamp(f, 0.0f, 1.0f) } {} - bool scatter(const ray& r_in, const hit_record& rec, - color& attenuation, ray& scattered) const { + bool scatter(const ray& r_in, const hit_record& rec, color& attenuation, + ray& scattered) const { auto rng = LocalPseudoRNG { toseed(rec.normal) }; vec reflected = reflect(unit_vector(r_in.direction()), rec.normal); scattered = ray(rec.p, reflected + fuzz * rng.in_unit_ball(), r_in.time()); @@ -66,8 +64,8 @@ struct dielectric_material { return r0 + (1 - r0) * sycl::pow((1 - cosine), 5.0f); } - bool scatter(const ray& r_in, const hit_record& rec, - color& attenuation, ray& scattered) const { + bool scatter(const ray& r_in, const hit_record& rec, color& attenuation, + ray& scattered) const { // Attenuation of the ray hitting the object is modified based on the color // at hit point auto rng = LocalPseudoRNG { toseed(rec.normal) }; @@ -130,6 +128,6 @@ struct isotropic_material { }; using material_t = - std::variant; + std::variant; #endif diff --git a/include/render.hpp b/include/render.hpp index cf2787c..1315bf7 100644 --- a/include/render.hpp +++ b/include/render.hpp @@ -14,8 +14,9 @@ #include "vec.hpp" #include "visit.hpp" -inline auto render_pixel(int width, int height, int depth, int samples, int x_coord, int y_coord, - camera const& cam, auto& hittable_acc, auto fb_acc) { +inline auto render_pixel(int width, int height, int depth, int samples, + int x_coord, int y_coord, camera const& cam, + auto& hittable_acc, auto fb_acc) { auto get_color = [&](const ray& r) { auto hit_world = [&](const ray& r, hit_record& rec, material_t& material_type) { @@ -97,14 +98,15 @@ inline auto render_pixel(int width, int height, int depth, int samples, int x_co struct PixelRender; -inline void executor(int width, int height, int depth, int samples, sycl::handler& cgh, - camera const& cam_ptr, auto& hittable_acc, auto& fb_acc) { +inline void executor(int width, int height, int depth, int samples, + sycl::handler& cgh, camera const& cam_ptr, + auto& hittable_acc, auto& fb_acc) { if constexpr (buildparams::use_single_task) { cgh.single_task([=] { for (int x_coord = 0; x_coord != width; ++x_coord) for (int y_coord = 0; y_coord != height; ++y_coord) { - render_pixel(width, height, depth, samples, x_coord, y_coord, - cam_ptr, hittable_acc, fb_acc); + render_pixel(width, height, depth, samples, x_coord, y_coord, cam_ptr, + hittable_acc, fb_acc); } }); } else { @@ -114,8 +116,8 @@ inline void executor(int width, int height, int depth, int samples, sycl::handle auto gid = item.get_id(); const auto x_coord = gid[1]; const auto y_coord = gid[0]; - render_pixel(width, height, depth, samples, x_coord, y_coord, - cam_ptr, hittable_acc, fb_acc); + render_pixel(width, height, depth, samples, x_coord, y_coord, cam_ptr, + hittable_acc, fb_acc); }); } } @@ -133,7 +135,6 @@ void render(int width, int height, int depth, int samples, sycl::queue& queue, auto fb_acc = frame_buf.get_access(cgh); auto hittables_acc = hittables_buf.get_access(cgh); - executor(width, height, depth, samples, cgh, cam, hittables_acc, - fb_acc); + executor(width, height, depth, samples, cgh, cam, hittables_acc, fb_acc); }); } diff --git a/include/rtweekend.hpp b/include/rtweekend.hpp index fb14e6a..22ee7a6 100644 --- a/include/rtweekend.hpp +++ b/include/rtweekend.hpp @@ -33,14 +33,16 @@ using vec = real_vec; // Utility Functions -inline real_t degrees_to_radians(real_t degrees) { return degrees * pi / 180.0f; } +inline real_t degrees_to_radians(real_t degrees) { + return degrees * pi / 180.0f; +} -uint32_t toseed(vec const & val) { +uint32_t toseed(vec const& val) { uint32_t x, y, z; std::memcpy(&x, &val.x(), sizeof(uint32_t)); std::memcpy(&y, &val.y(), sizeof(uint32_t)); std::memcpy(&z, &val.z(), sizeof(uint32_t)); - return x*x*y*y*z*z; + return x * x * y * y * z * z; } class LocalPseudoRNG { @@ -55,7 +57,7 @@ class LocalPseudoRNG { } // Returns a random float in min, max - inline real_t real(real_t min, real_t max) { + inline real_t real(real_t min, real_t max) { // TODO use FMA ? return min + (max - min) * real(); } diff --git a/include/visit.hpp b/include/visit.hpp index 23fde8f..b1417ee 100644 --- a/include/visit.hpp +++ b/include/visit.hpp @@ -47,7 +47,8 @@ decltype(auto) visit_single(Func&& f, Var&& var) { auto monostate_dispatch(auto&& dispatch, auto&& monostate_value) { return [&](auto&& arg) { - if constexpr (std::is_same_v, std::monostate>) { + if constexpr (std::is_same_v, + std::monostate>) { assert(false && "Try to dispatch to monostate value"); return monostate_value; } else { From 70c34f9273c497498eb9339904acdaae65254c8b Mon Sep 17 00:00:00 2001 From: Luc Forget Date: Fri, 11 Jun 2021 00:58:58 -0700 Subject: [PATCH 10/14] Improving decentralised RNG seed --- include/constant_medium.hpp | 3 ++- include/material.hpp | 8 ++++---- 2 files changed, 6 insertions(+), 5 deletions(-) diff --git a/include/constant_medium.hpp b/include/constant_medium.hpp index 8ee509b..4d9cce5 100644 --- a/include/constant_medium.hpp +++ b/include/constant_medium.hpp @@ -3,6 +3,7 @@ #include "box.hpp" #include "material.hpp" +#include "rtweekend.hpp" #include "sphere.hpp" #include "texture.hpp" #include "visit.hpp" @@ -59,7 +60,7 @@ class constant_medium { /// Distance between the two hitpoints affect of probability /// of the ray hitting a smoke particle const auto distance_inside_boundary = (rec2.t - rec1.t) * ray_length; - auto rng = LocalPseudoRNG { toseed(r.direction()) }; + auto rng = LocalPseudoRNG { toseed(r.direction()) ^ toseed(r.origin()) }; const auto hit_distance = neg_inv_density * sycl::log(rng.real()); /// With lower density, hit_distance has higher probabilty diff --git a/include/material.hpp b/include/material.hpp index 0dfc648..ff642fe 100644 --- a/include/material.hpp +++ b/include/material.hpp @@ -17,7 +17,7 @@ struct lambertian_material { bool scatter(const ray& r_in, const hit_record& rec, color& attenuation, ray& scattered) const { - auto rng = LocalPseudoRNG { toseed(rec.normal) }; + LocalPseudoRNG rng { toseed(r_in.direction()) ^ toseed(r_in.origin()) }; vec scatter_direction = rec.normal + rng.unit_vec(); scattered = ray(rec.p, scatter_direction, r_in.time()); // Attenuation of the ray hitting the object is modified based on the color @@ -37,7 +37,7 @@ struct metal_material { bool scatter(const ray& r_in, const hit_record& rec, color& attenuation, ray& scattered) const { - auto rng = LocalPseudoRNG { toseed(rec.normal) }; + LocalPseudoRNG rng { toseed(r_in.direction()) ^ toseed(r_in.origin()) }; vec reflected = reflect(unit_vector(r_in.direction()), rec.normal); scattered = ray(rec.p, reflected + fuzz * rng.in_unit_ball(), r_in.time()); // Attenuation of the ray hitting the object is modified based on the color @@ -68,7 +68,7 @@ struct dielectric_material { ray& scattered) const { // Attenuation of the ray hitting the object is modified based on the color // at hit point - auto rng = LocalPseudoRNG { toseed(rec.normal) }; + LocalPseudoRNG rng { toseed(r_in.direction()) ^ toseed(r_in.origin()) }; attenuation *= albedo; real_t refraction_ratio = rec.front_face ? (1.0f / ref_idx) : ref_idx; vec unit_direction = unit_vector(r_in.direction()); @@ -116,7 +116,7 @@ struct isotropic_material { bool scatter(const ray& r_in, const hit_record& rec, color& attenuation, ray& scattered) const { - LocalPseudoRNG rng(toseed(r_in.direction())); + LocalPseudoRNG rng { toseed(r_in.direction()) ^ toseed(r_in.origin()) }; scattered = ray(rec.p, rng.in_unit_ball(), r_in.time()); attenuation *= dev_visit([&](auto&& t) { return t.value(rec); }, albedo); return true; From b92d9c279f9e3be430f182c5d1f3bb714b488b69 Mon Sep 17 00:00:00 2001 From: Luc Forget Date: Fri, 11 Jun 2021 06:08:20 -0700 Subject: [PATCH 11/14] FPGA friendly local random seed generator --- include/constant_medium.hpp | 2 +- include/material.hpp | 8 ++++---- include/rtweekend.hpp | 18 ++++++++++++++++++ src/main.cpp | 30 ++++++++++++++++++++---------- 4 files changed, 43 insertions(+), 15 deletions(-) diff --git a/include/constant_medium.hpp b/include/constant_medium.hpp index 4d9cce5..42ac675 100644 --- a/include/constant_medium.hpp +++ b/include/constant_medium.hpp @@ -60,7 +60,7 @@ class constant_medium { /// Distance between the two hitpoints affect of probability /// of the ray hitting a smoke particle const auto distance_inside_boundary = (rec2.t - rec1.t) * ray_length; - auto rng = LocalPseudoRNG { toseed(r.direction()) ^ toseed(r.origin()) }; + auto rng = LocalPseudoRNG { toseed(r.direction(), r.origin()) }; const auto hit_distance = neg_inv_density * sycl::log(rng.real()); /// With lower density, hit_distance has higher probabilty diff --git a/include/material.hpp b/include/material.hpp index ff642fe..bee4238 100644 --- a/include/material.hpp +++ b/include/material.hpp @@ -17,7 +17,7 @@ struct lambertian_material { bool scatter(const ray& r_in, const hit_record& rec, color& attenuation, ray& scattered) const { - LocalPseudoRNG rng { toseed(r_in.direction()) ^ toseed(r_in.origin()) }; + LocalPseudoRNG rng { toseed(r_in.direction(), r_in.origin()) }; vec scatter_direction = rec.normal + rng.unit_vec(); scattered = ray(rec.p, scatter_direction, r_in.time()); // Attenuation of the ray hitting the object is modified based on the color @@ -37,7 +37,7 @@ struct metal_material { bool scatter(const ray& r_in, const hit_record& rec, color& attenuation, ray& scattered) const { - LocalPseudoRNG rng { toseed(r_in.direction()) ^ toseed(r_in.origin()) }; + LocalPseudoRNG rng { toseed(r_in.direction(), r_in.origin()) }; vec reflected = reflect(unit_vector(r_in.direction()), rec.normal); scattered = ray(rec.p, reflected + fuzz * rng.in_unit_ball(), r_in.time()); // Attenuation of the ray hitting the object is modified based on the color @@ -68,7 +68,7 @@ struct dielectric_material { ray& scattered) const { // Attenuation of the ray hitting the object is modified based on the color // at hit point - LocalPseudoRNG rng { toseed(r_in.direction()) ^ toseed(r_in.origin()) }; + LocalPseudoRNG rng { toseed(r_in.direction(), r_in.origin()) }; attenuation *= albedo; real_t refraction_ratio = rec.front_face ? (1.0f / ref_idx) : ref_idx; vec unit_direction = unit_vector(r_in.direction()); @@ -116,7 +116,7 @@ struct isotropic_material { bool scatter(const ray& r_in, const hit_record& rec, color& attenuation, ray& scattered) const { - LocalPseudoRNG rng { toseed(r_in.direction()) ^ toseed(r_in.origin()) }; + LocalPseudoRNG rng { toseed(r_in.direction(), r_in.origin()) }; scattered = ray(rec.p, rng.in_unit_ball(), r_in.time()); attenuation *= dev_visit([&](auto&& t) { return t.value(rec); }, albedo); return true; diff --git a/include/rtweekend.hpp b/include/rtweekend.hpp index 22ee7a6..514831e 100644 --- a/include/rtweekend.hpp +++ b/include/rtweekend.hpp @@ -3,6 +3,7 @@ #include #include +#include #include #include #include @@ -45,6 +46,23 @@ uint32_t toseed(vec const& val) { return x * x * y * y * z * z; } +uint32_t toseed(vec const& val1, vec const& val2) { + uint32_t x1, y1, z1, x2, y2, z2; + std::memcpy(&x1, &val1.x(), sizeof(uint32_t)); + std::memcpy(&y1, &val1.y(), sizeof(uint32_t)); + std::memcpy(&z1, &val1.z(), sizeof(uint32_t)); + std::memcpy(&x2, &val2.x(), sizeof(uint32_t)); + std::memcpy(&y2, &val2.y(), sizeof(uint32_t)); + std::memcpy(&z2, &val2.z(), sizeof(uint32_t)); + uint32_t shifted1 = x1 << 26; + uint32_t shifted2 = (x2 & 63) << 21; + uint32_t shifted3 = (y1 & 63) << 15; + uint32_t shifted4 = (y2 & 63) << 10; + uint32_t shifted5 = (z1 & 63) << 5; + uint32_t shifted6 = (z2 & 63); + return shifted1 ^ shifted2 ^ shifted3 ^ shifted4 ^ shifted5 ^ shifted6; +} + class LocalPseudoRNG { public: inline LocalPseudoRNG(std::uint32_t init_state = xorshift<>::initial_state) diff --git a/src/main.cpp b/src/main.cpp index dea2ea4..2bc1697 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -2,6 +2,7 @@ #include #include #include +#include #include #include #include @@ -79,17 +80,29 @@ int main(int argc, char* argv[]) { hittables.emplace_back(sphere(point { 0, -1000, 0 }, 1000, m)); t = checker_texture(color { 0.9f, 0.9f, 0.9f }, color { 0.4f, 0.2f, 0.1f }); - LocalPseudoRNG rng; + LocalPseudoRNG rng{static_cast(std::time(nullptr))}; - for (int a = -11; a < 11; a += 3) { - for (int b = -11; b < 11; b += 3) { + for (int a = -11; a < 11; a++) { + for (int b = -11; b < 11; b++) { + auto choose_mat = rng.real(); // Spheres are placed at a point randomly displaced from a,b point center(a + 0.9f * rng.real(), 0.2f, b + 0.9f * rng.real()); - if (sycl::length((center - point(4, 0.2f, 0))) > 0.9f) { + if (choose_mat < 0.70f) { // Lambertian auto albedo = rng.vec_t() * rng.vec_t(); hittables.emplace_back( sphere(center, 0.2f, lambertian_material(albedo))); + } else if (choose_mat < 0.95f) { + // metal + auto albedo = rng.vec_t(0.5f, 1); + auto fuzz = rng.real(0, 0.5f); + hittables.emplace_back( + sphere(center, 0.2f, metal_material(albedo, fuzz))); + } else { + // glass + hittables.emplace_back( + sphere(center, 0.2f, + dielectric_material(1.5f, color { 1.0f, 1.0f, 1.0f }))); } } } @@ -127,23 +140,20 @@ int main(int argc, char* argv[]) { hittables.emplace_back(sphere(point { 0, 1, -2.25f }, 1, metal_material(color(0.7f, 0.6f, 0.5f), 0.0f))); - // t = image_texture::image_texture_factory("../images/SYCL.png", 5); - - // // Add a sphere with a SYCL logo in the background hittables.emplace_back( sphere { point { -60, 3, 5 }, 4, lambertian_material { t } }); // Add a metallic monolith hittables.emplace_back( box { point { 6.5f, 0, -1.5f }, point { 7.0f, 3.0f, -1.0f }, - metal_material { color { 0.7f, 0.6f, 0.5f }, 0.25f } });/**/ + metal_material { color { 0.7f, 0.6f, 0.5f }, 0.25f } }); // Add a smoke ball sphere smoke_sphere = sphere { point { 5, 1, 3.5f }, 1, lambertian_material { color { 0.75f, 0.75f, 0.75f } } }; - /*hittables.emplace_back( - constant_medium { smoke_sphere, 1, color { 1, 1, 1 } });*/ + hittables.emplace_back( + constant_medium { smoke_sphere, 1, color { 1, 1, 1 } }); // SYCL queue sycl::queue myQueue; From 635037d4cf3b71c324eb2aa1f66bd5715bb048eb Mon Sep 17 00:00:00 2001 From: Luc Forget Date: Tue, 15 Jun 2021 01:27:26 -0700 Subject: [PATCH 12/14] Address review comments --- README.md | 11 +++++++++-- include/rtweekend.hpp | 23 +++++++++++------------ src/main.cpp | 21 ++++++++++++++++----- 3 files changed, 36 insertions(+), 19 deletions(-) diff --git a/README.md b/README.md index cb4a1ee..873ca17 100644 --- a/README.md +++ b/README.md @@ -91,9 +91,16 @@ This creates the executable. Now you can run the path tracer with: ```sh -time RT_SYCL/sycl-rt >! result.ppm +time ./sycl-rt 800 480 50 100 ``` -This results in the image ``result.ppm`` produced by the path tracer. +This results in the image ``out.png`` produced by the path tracer. + +Parameters to the executable are + ++ The output image width (here 800) ++ The output image height (here 480) ++ The maximum bouncing depth of the ray (here 50) ++ The number of samples per pixel (here 100) ## Bibliography diff --git a/include/rtweekend.hpp b/include/rtweekend.hpp index 514831e..99c46f6 100644 --- a/include/rtweekend.hpp +++ b/include/rtweekend.hpp @@ -38,14 +38,13 @@ inline real_t degrees_to_radians(real_t degrees) { return degrees * pi / 180.0f; } -uint32_t toseed(vec const& val) { - uint32_t x, y, z; - std::memcpy(&x, &val.x(), sizeof(uint32_t)); - std::memcpy(&y, &val.y(), sizeof(uint32_t)); - std::memcpy(&z, &val.z(), sizeof(uint32_t)); - return x * x * y * y * z * z; -} +/** + @brief hash two vector using coordinates low bits + @param val1 first vector + @param val2 second vector + @return uint32_t + */ uint32_t toseed(vec const& val1, vec const& val2) { uint32_t x1, y1, z1, x2, y2, z2; std::memcpy(&x1, &val1.x(), sizeof(uint32_t)); @@ -55,12 +54,12 @@ uint32_t toseed(vec const& val1, vec const& val2) { std::memcpy(&y2, &val2.y(), sizeof(uint32_t)); std::memcpy(&z2, &val2.z(), sizeof(uint32_t)); uint32_t shifted1 = x1 << 26; - uint32_t shifted2 = (x2 & 63) << 21; + uint32_t shifted2 = (x2 & 31) << 21; uint32_t shifted3 = (y1 & 63) << 15; - uint32_t shifted4 = (y2 & 63) << 10; - uint32_t shifted5 = (z1 & 63) << 5; - uint32_t shifted6 = (z2 & 63); - return shifted1 ^ shifted2 ^ shifted3 ^ shifted4 ^ shifted5 ^ shifted6; + uint32_t shifted4 = (y2 & 31) << 10; + uint32_t shifted5 = (z1 & 31) << 5; + uint32_t shifted6 = (z2 & 31); + return shifted1 | shifted2 | shifted3 | shifted4 | shifted5 | shifted6; } class LocalPseudoRNG { diff --git a/src/main.cpp b/src/main.cpp index 2bc1697..4441081 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -1,4 +1,5 @@ #include "sycl.hpp" +#include "xorshift.hpp" #include #include #include @@ -6,6 +7,7 @@ #include #include #include +#include #include #include @@ -59,8 +61,9 @@ void save_image_png(int width, int height, sycl::buffer& fb) { } int main(int argc, char* argv[]) { - if (argc != 5) { - std::cerr << "Usage: sycl-rt OUT_WIDTH OUT_HEIGHT DEPTH SAMPLES" + if (argc < 5 || argc > 7) { + std::cerr << "Usage: sycl-rt OUT_WIDTH OUT_HEIGHT DEPTH SAMPLES " + "[SPHERE_INC [RAND_SEED]]" << std::endl; return -1; } @@ -69,6 +72,14 @@ int main(int argc, char* argv[]) { auto height = std::stoi({ argv[2] }); auto depth = std::stoi({ argv[3] }); auto samples = std::stoi({ argv[4] }); + int sphere_inc = 1; + if (argc >= 6) + sphere_inc = std::stoi({ argv[5] }); + + auto rand_seed = xorshift<>::initial_state; + + if (argc >= 7) + rand_seed = std::stoi({ argv[6] }); /// Graphical objects std::vector hittables; @@ -80,10 +91,10 @@ int main(int argc, char* argv[]) { hittables.emplace_back(sphere(point { 0, -1000, 0 }, 1000, m)); t = checker_texture(color { 0.9f, 0.9f, 0.9f }, color { 0.4f, 0.2f, 0.1f }); - LocalPseudoRNG rng{static_cast(std::time(nullptr))}; + LocalPseudoRNG rng{rand_seed}; - for (int a = -11; a < 11; a++) { - for (int b = -11; b < 11; b++) { + for (int a = -11; a < 11; a += sphere_inc) { + for (int b = -11; b < 11; b += sphere_inc) { auto choose_mat = rng.real(); // Spheres are placed at a point randomly displaced from a,b point center(a + 0.9f * rng.real(), 0.2f, b + 0.9f * rng.real()); From 5e968175cc63e7e81009be5c08a42d740793e3f1 Mon Sep 17 00:00:00 2001 From: Luc Forget Date: Tue, 15 Jun 2021 04:38:18 -0700 Subject: [PATCH 13/14] Restore motion blur --- include/camera.hpp | 2 +- include/render.hpp | 7 +++++-- src/main.cpp | 45 +++++++++++++++++++++++++++------------------ 3 files changed, 33 insertions(+), 21 deletions(-) diff --git a/include/camera.hpp b/include/camera.hpp index f67ddd5..8265bcf 100644 --- a/include/camera.hpp +++ b/include/camera.hpp @@ -96,7 +96,7 @@ class camera { return { origin + offset, lower_left_corner + s * horizontal + t * vertical - origin - offset, - time0 }; + rng.real(time0, time1) }; } }; diff --git a/include/render.hpp b/include/render.hpp index 1315bf7..b2122e7 100644 --- a/include/render.hpp +++ b/include/render.hpp @@ -80,13 +80,16 @@ inline auto render_pixel(int width, int height, int depth, int samples, // If not returned within max_depth return black return color { 0.0f, 0.0f, 0.0f }; }; - uint32_t seed = std::hash {}((x_coord << 16) | (y_coord & 0xFFFF)); - LocalPseudoRNG rng(seed); + color final_color(0.0f, 0.0f, 0.0f); + + uint32_t seed = ~(x_coord << 16 | (y_coord & 0xFFFF)); + LocalPseudoRNG rng { seed }; for (auto i = 0; i < samples; i++) { const auto u = (x_coord + rng.real()) / width; const auto v = (y_coord + rng.real()) / height; // u and v are points on the viewport + ray r = cam.get_ray(u, v, rng); final_color += get_color(r); } diff --git a/src/main.cpp b/src/main.cpp index 4441081..1af66e4 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -91,29 +91,38 @@ int main(int argc, char* argv[]) { hittables.emplace_back(sphere(point { 0, -1000, 0 }, 1000, m)); t = checker_texture(color { 0.9f, 0.9f, 0.9f }, color { 0.4f, 0.2f, 0.1f }); - LocalPseudoRNG rng{rand_seed}; + LocalPseudoRNG rng { rand_seed }; for (int a = -11; a < 11; a += sphere_inc) { for (int b = -11; b < 11; b += sphere_inc) { + // Based on a random variable , the material type is chosen auto choose_mat = rng.real(); // Spheres are placed at a point randomly displaced from a,b point center(a + 0.9f * rng.real(), 0.2f, b + 0.9f * rng.real()); - if (choose_mat < 0.70f) { - // Lambertian - auto albedo = rng.vec_t() * rng.vec_t(); - hittables.emplace_back( - sphere(center, 0.2f, lambertian_material(albedo))); - } else if (choose_mat < 0.95f) { - // metal - auto albedo = rng.vec_t(0.5f, 1); - auto fuzz = rng.real(0, 0.5f); - hittables.emplace_back( - sphere(center, 0.2f, metal_material(albedo, fuzz))); - } else { - // glass - hittables.emplace_back( - sphere(center, 0.2f, - dielectric_material(1.5f, color { 1.0f, 1.0f, 1.0f }))); + if (sycl::length((center - point(4, 0.2f, 0))) > 0.9f) { + if (choose_mat < 0.4f) { + // Lambertian + auto albedo = rng.vec_t() * rng.vec_t(); + hittables.emplace_back( + sphere(center, 0.2f, lambertian_material(albedo))); + } else if (choose_mat < 0.8f) { + // Lambertian movig spheres + auto albedo = rng.vec_t() * rng.vec_t(); + auto center2 = center + point { 0, rng.real(0, 0.25f), 0 }; + hittables.emplace_back(sphere(center, center2, 0.0f, 1.0f, 0.2f, + lambertian_material(albedo))); + } else if (choose_mat < 0.95f) { + // metal + auto albedo = rng.vec_t(0.5f, 1); + auto fuzz = rng.real(0, 0.5f); + hittables.emplace_back( + sphere(center, 0.2f, metal_material(albedo, fuzz))); + } else { + // glass + hittables.emplace_back( + sphere(center, 0.2f, + dielectric_material(1.5f, color { 1.0f, 1.0f, 1.0f }))); + } } } } @@ -140,7 +149,7 @@ int main(int argc, char* argv[]) { // Four large spheres of metal, dielectric and Lambertian material types // t = image_texture::image_texture_factory("../images/Xilinx.jpg"); - //hittables.emplace_back(xy_rect(2, 4, 0, 1, -1, lambertian_material(t))); + // hittables.emplace_back(xy_rect(2, 4, 0, 1, -1, lambertian_material(t))); hittables.emplace_back( sphere(point { 4, 1, 2.25f }, 1, lambertian_material(t))); hittables.emplace_back( From 092b972e54d56867979cda458a2e6333e42d710e Mon Sep 17 00:00:00 2001 From: Luc Forget Date: Tue, 15 Jun 2021 05:23:44 -0700 Subject: [PATCH 14/14] Restore image texture and xy_rect --- include/hitable.hpp | 2 +- include/material.hpp | 24 +++++----- include/render.hpp | 20 ++++---- include/rtweekend.hpp | 3 +- include/texture.hpp | 105 ++++++++++++++++++++++++++++++++++++++++-- src/main.cpp | 10 ++-- 6 files changed, 134 insertions(+), 30 deletions(-) diff --git a/include/hitable.hpp b/include/hitable.hpp index 41a8ff1..e7e0462 100644 --- a/include/hitable.hpp +++ b/include/hitable.hpp @@ -7,5 +7,5 @@ #include "sphere.hpp" #include "triangle.hpp" -using hittable_t = std::variant; +using hittable_t = std::variant; #endif diff --git a/include/material.hpp b/include/material.hpp index bee4238..46c9486 100644 --- a/include/material.hpp +++ b/include/material.hpp @@ -15,17 +15,17 @@ struct lambertian_material { lambertian_material(const texture_t& a) : albedo { a } {} - bool scatter(const ray& r_in, const hit_record& rec, color& attenuation, + bool scatter(auto& ctx, const ray& r_in, const hit_record& rec, color& attenuation, ray& scattered) const { LocalPseudoRNG rng { toseed(r_in.direction(), r_in.origin()) }; vec scatter_direction = rec.normal + rng.unit_vec(); scattered = ray(rec.p, scatter_direction, r_in.time()); // Attenuation of the ray hitting the object is modified based on the color // at hit point - attenuation *= dev_visit([&](auto&& t) { return t.value(rec); }, albedo); + attenuation *= dev_visit([&](auto&& t) { return t.value(ctx, rec); }, albedo); return true; } - color emitted(const hit_record& rec) { return color(0.f, 0.f, 0.f); } + color emitted(auto&, const hit_record& rec) { return color(0.f, 0.f, 0.f); } texture_t albedo; }; @@ -35,7 +35,7 @@ struct metal_material { : albedo { a } , fuzz { std::clamp(f, 0.0f, 1.0f) } {} - bool scatter(const ray& r_in, const hit_record& rec, color& attenuation, + bool scatter(auto&, const ray& r_in, const hit_record& rec, color& attenuation, ray& scattered) const { LocalPseudoRNG rng { toseed(r_in.direction(), r_in.origin()) }; vec reflected = reflect(unit_vector(r_in.direction()), rec.normal); @@ -46,7 +46,7 @@ struct metal_material { return (dot(scattered.direction(), rec.normal) > 0); } - color emitted(const hit_record& rec) { return color(0, 0, 0); } + color emitted(auto&, const hit_record& rec) { return color(0, 0, 0); } color albedo; real_t fuzz; }; @@ -64,7 +64,7 @@ struct dielectric_material { return r0 + (1 - r0) * sycl::pow((1 - cosine), 5.0f); } - bool scatter(const ray& r_in, const hit_record& rec, color& attenuation, + bool scatter(auto&, const ray& r_in, const hit_record& rec, color& attenuation, ray& scattered) const { // Attenuation of the ray hitting the object is modified based on the color // at hit point @@ -85,7 +85,7 @@ struct dielectric_material { return true; } - color emitted(const hit_record& rec) { return color(0, 0, 0); } + color emitted(auto&, const hit_record& rec) { return color(0, 0, 0); } // Refractive index of the glass real_t ref_idx; // Color of the glass @@ -101,8 +101,8 @@ struct lightsource_material { template bool scatter(T&...) const { return false; } - color emitted(const hit_record& rec) { - return dev_visit([&](auto&& t) { return t.value(rec); }, emit); + color emitted(auto& ctx, const hit_record& rec) { + return dev_visit([&](auto&& t) { return t.value(ctx, rec); }, emit); } texture_t emit; @@ -114,15 +114,15 @@ struct isotropic_material { isotropic_material(texture_t& a) : albedo { a } {} - bool scatter(const ray& r_in, const hit_record& rec, color& attenuation, + bool scatter(auto& ctx, const ray& r_in, const hit_record& rec, color& attenuation, ray& scattered) const { LocalPseudoRNG rng { toseed(r_in.direction(), r_in.origin()) }; scattered = ray(rec.p, rng.in_unit_ball(), r_in.time()); - attenuation *= dev_visit([&](auto&& t) { return t.value(rec); }, albedo); + attenuation *= dev_visit([&](auto&& t) { return t.value(ctx, rec); }, albedo); return true; } - color emitted(const hit_record& rec) { return color(0, 0, 0); } + color emitted(auto&, const hit_record& rec) { return color(0, 0, 0); } texture_t albedo; }; diff --git a/include/render.hpp b/include/render.hpp index b2122e7..91efdb6 100644 --- a/include/render.hpp +++ b/include/render.hpp @@ -14,7 +14,7 @@ #include "vec.hpp" #include "visit.hpp" -inline auto render_pixel(int width, int height, int depth, int samples, +inline auto render_pixel(auto& ctx, int width, int height, int depth, int samples, int x_coord, int y_coord, camera const& cam, auto& hittable_acc, auto fb_acc) { auto get_color = [&](const ray& r) { @@ -49,11 +49,11 @@ inline auto render_pixel(int width, int height, int depth, int samples, for (auto i = 0; i < depth; i++) { hit_record rec; if (hit_world(cur_ray, rec, material_type)) { - emitted = dev_visit([&](auto&& mat) { return mat.emitted(rec); }, + emitted = dev_visit([&](auto&& mat) { return mat.emitted(ctx, rec); }, material_type); if (dev_visit( [&](auto&& mat) { - return mat.scatter(cur_ray, rec, cur_attenuation, scattered); + return mat.scatter(ctx, cur_ray, rec, cur_attenuation, scattered); }, material_type)) { // On hitting the object, the ray gets scattered @@ -103,13 +103,14 @@ struct PixelRender; inline void executor(int width, int height, int depth, int samples, sycl::handler& cgh, camera const& cam_ptr, - auto& hittable_acc, auto& fb_acc) { + auto& hittable_acc, auto& fb_acc, auto& texture_acc) { if constexpr (buildparams::use_single_task) { cgh.single_task([=] { + task_context ctx { texture_acc.get_pointer() }; for (int x_coord = 0; x_coord != width; ++x_coord) for (int y_coord = 0; y_coord != height; ++y_coord) { - render_pixel(width, height, depth, samples, x_coord, y_coord, cam_ptr, - hittable_acc, fb_acc); + render_pixel(ctx, width, height, depth, samples, x_coord, y_coord, + cam_ptr, hittable_acc, fb_acc); } }); } else { @@ -119,7 +120,8 @@ inline void executor(int width, int height, int depth, int samples, auto gid = item.get_id(); const auto x_coord = gid[1]; const auto y_coord = gid[0]; - render_pixel(width, height, depth, samples, x_coord, y_coord, cam_ptr, + task_context ctx { texture_acc.get_pointer() }; + render_pixel(ctx, width, height, depth, samples, x_coord, y_coord, cam_ptr, hittable_acc, fb_acc); }); } @@ -132,12 +134,14 @@ void render(int width, int height, int depth, int samples, sycl::queue& queue, const auto nb_hittable = hittables.size(); auto hittables_buf = sycl::buffer(hittables.data(), sycl::range<1>(nb_hittable)); + auto texture_buf = image_texture::freeze(); // Submit command group on device queue.submit([&](sycl::handler& cgh) { auto fb_acc = frame_buf.get_access(cgh); auto hittables_acc = hittables_buf.get_access(cgh); - executor(width, height, depth, samples, cgh, cam, hittables_acc, fb_acc); + auto texture_acc = texture_buf.get_access(cgh); + executor(width, height, depth, samples, cgh, cam, hittables_acc, fb_acc, texture_acc); }); } diff --git a/include/rtweekend.hpp b/include/rtweekend.hpp index 99c46f6..af42e2c 100644 --- a/include/rtweekend.hpp +++ b/include/rtweekend.hpp @@ -129,7 +129,8 @@ class LocalPseudoRNG { kernel callees */ struct task_context { - LocalPseudoRNG rng; + // See image_texture in texture.hpp for more details + sycl::global_ptr texture_data; }; // Common Headers diff --git a/include/texture.hpp b/include/texture.hpp index cece438..e0a42b0 100644 --- a/include/texture.hpp +++ b/include/texture.hpp @@ -22,7 +22,7 @@ struct solid_texture { solid_texture(real_t red, real_t green, real_t blue) : solid_texture { color { red, green, blue } } {} // For solid texture, the color is same throughout the sphere - color value(const hit_record&) const { return color_value; } + color value(auto&, const hit_record&) const { return color_value; } private: color color_value; @@ -39,18 +39,113 @@ struct checker_texture { : odd { solid_texture { c1 } } , even { solid_texture { c2 } } {} // Color value is different based on normalised spherical coordinates - color value(const hit_record& rec) const { + color value(auto& ctx, const hit_record& rec) const { auto sines = sycl::sin(10 * rec.p.x()) * sycl::sin(10 * rec.p.y()) * sycl::sin(10 * rec.p.z()); if (sines < 0) - return odd.value(rec); + return odd.value(ctx, rec); else - return even.value(rec); + return even.value(ctx, rec); } solid_texture odd; solid_texture even; }; -using texture_t = std::variant; +/** + @brief A texture based on an image + In order to be able to get the bitmap on the device without embedding it in + the object, all image_texture textures are serialized in one vector. + The offset of the texture in the vector is stored in the image_texture + instance. + When all the textures have been loaded, the freeze() method can be called to + get a sycl::buffer that store this data. + */ +struct image_texture { + private: + static constexpr auto bytes_per_pixel = 3; + // Vector in which all the textures are serialized + static std::vector texture_data; + static bool frozen; + + std::size_t width {}; + std::size_t height {}; + // offset of the first pixel in the texture vector + std::size_t offset; + + /// The repetition rate of the image + float cyclic_frequency { 1.f }; + + image_texture(std::size_t _width, std::size_t _height, std::size_t _offset, + float _cyclic_frequency) + : width { _width } + , height { _height } + , offset { _offset } + , cyclic_frequency { _cyclic_frequency } {} + + public: + /** Create a texture from an image file + \param[in] file_name is the path name to the image file + \param[in] cyclic_frequency is an optional repetition rate of + the image in the texture + */ + static image_texture image_texture_factory(const char* file_name, + float _cyclic_frequency = 1) { + assert(!frozen); + auto components_per_pixel = bytes_per_pixel; + uint8_t* _data; + int _w, _h; + _data = + stbi_load(file_name, &_w, &_h, &components_per_pixel, bytes_per_pixel); + std::size_t _offset = 0; + if (!_data) { + std::cerr << "ERROR: Could not load texture image file '" << file_name + << "'.\n" + << stbi_failure_reason() << std::endl; + _w = _h = 1; + } else { + auto size = bytes_per_pixel * _w * _h; + _offset = texture_data.size() / 3; + std::copy(_data, _data + size, std::back_inserter(texture_data)); + } + return image_texture(_w, _h, _offset, _cyclic_frequency); + } + + /** + @brief Get a sycl::buffer containing texture data. + image_texture_factory should not be called after having called freeze + @return sycl::buffer + */ + static sycl::buffer freeze() { + assert(!frozen); + frozen = true; + return sycl::buffer { texture_data.data(), + { texture_data.size() / 3, 3 } }; + } + + /// Get the color for the texture at the given place + /// \todo rename this value() to color() everywhere? + color value(auto& ctx, const hit_record& rec) const { + // If texture data is unavailable, return solid cyan + // The image is repeated by the repetition factor + + std::size_t i = + sycl::fmod(rec.u * cyclic_frequency, (float)1) * (width - 1); + // The image frame buffer is going downwards, so flip the y axis + std::size_t j = + (1 - sycl::fmod(rec.v * cyclic_frequency, (float)1)) * (height - 1); + std::size_t local_offset = j * width + i; + std::size_t pix_idx = local_offset + offset; + auto scale = 1.f / 255; + auto& texture_data = ctx.texture_data; + return { texture_data[pix_idx * 3] * scale, + texture_data[pix_idx * 3 + 1] * scale, + texture_data[pix_idx * 3 + 2] * scale }; + } +}; + +using texture_t = std::variant; +// Start filled with the fallback texture (solid blue) for texture load error +std::vector image_texture::texture_data { 0, 0, 1 }; +bool image_texture::frozen = false; #endif diff --git a/src/main.cpp b/src/main.cpp index 1af66e4..e7b6bf2 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -147,11 +147,13 @@ int main(int argc, char* argv[]) { hittables.emplace_back( sphere(point { 4, 1, 0 }, 0.2f, lightsource_material(color(10, 0, 10)))); - // Four large spheres of metal, dielectric and Lambertian material types - // t = image_texture::image_texture_factory("../images/Xilinx.jpg"); - // hittables.emplace_back(xy_rect(2, 4, 0, 1, -1, lambertian_material(t))); + // Xilinx logo rectangle and sphere + t = image_texture::image_texture_factory("../images/Xilinx.jpg"); + hittables.emplace_back(xy_rect(2, 4, 0, 1, -1, lambertian_material(t))); hittables.emplace_back( sphere(point { 4, 1, 2.25f }, 1, lambertian_material(t))); + + // Four large spheres of metal, dielectric and Lambertian material types hittables.emplace_back( sphere(point { 0, 1, 0 }, 1, dielectric_material(1.5f, color { 1.0f, 0.5f, 0.5f }))); @@ -160,6 +162,8 @@ int main(int argc, char* argv[]) { hittables.emplace_back(sphere(point { 0, 1, -2.25f }, 1, metal_material(color(0.7f, 0.6f, 0.5f), 0.0f))); + //Add a sphere with a SYCL logo in the background + t = image_texture::image_texture_factory("../images/SYCL.png", 5); hittables.emplace_back( sphere { point { -60, 3, 5 }, 4, lambertian_material { t } });