diff --git a/CMakeLists.txt b/CMakeLists.txt index 92d0ae6..ffc030e 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -41,18 +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() - set(SYCL_RT_SRC_DIR ${CMAKE_CURRENT_SOURCE_DIR}/src) set(SYCL_RT_INCLUDE_DIR ${CMAKE_CURRENT_SOURCE_DIR}/include) @@ -60,8 +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}) # This is a SYCL program if ("${SYCL_CXX_COMPILER}" STREQUAL "") diff --git a/README.md b/README.md index 9e19d51..f76d2cb 100644 --- a/README.md +++ b/README.md @@ -93,9 +93,16 @@ This creates the executable. Now you can run the path tracer with: ```sh -time ./sycl-rt +time ./sycl-rt 800 480 50 100 ``` -This results in the image `out.png` 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/box.hpp b/include/box.hpp index feea59f..e0bcb82 100644 --- a/include/box.hpp +++ b/include/box.hpp @@ -26,7 +26,7 @@ class box { } /// Compute ray interaction with the box - bool hit(auto& ctx, 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_record temp_rec; material_t temp_material_type; @@ -36,7 +36,7 @@ class box { for (const auto& side : sides) { if (dev_visit( [&](auto&& arg) { - return arg.hit(ctx, r, min, closest_so_far, temp_rec, + return arg.hit(r, min, closest_so_far, temp_rec, temp_material_type); }, side)) { diff --git a/include/build_parameters.hpp b/include/build_parameters.hpp index 5ecc388..52054eb 100644 --- a/include/build_parameters.hpp +++ b/include/build_parameters.hpp @@ -14,8 +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 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..42ac675 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" @@ -25,24 +26,21 @@ class constant_medium { , 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, + bool hit(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( [&](auto&& arg) { - return arg.hit(ctx, r, -infinity, infinity, rec1, - temp_material_type); + return arg.hit(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, + return arg.hit(r, rec1.t + 0.0001f, infinity, rec2, temp_material_type); }, boundary)) { @@ -62,7 +60,8 @@ 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()); + 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 /// of being greater than distance_inside_boundary diff --git a/include/hit_record.hpp b/include/hit_record.hpp new file mode 100644 index 0000000..7b870dc --- /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 d5eb79a..e7e0462 100644 --- a/include/hitable.hpp +++ b/include/hitable.hpp @@ -1,26 +1,11 @@ #ifndef HITTABLE_H #define HITTABLE_H +#include "box.hpp" +#include "constant_medium.hpp" #include "ray.hpp" -#include "rtweekend.hpp" -#include "vec.hpp" - -class hit_record { - public: - float 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; - - // 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; - } -}; +#include "sphere.hpp" +#include "triangle.hpp" +using hittable_t = std::variant; #endif diff --git a/include/material.hpp b/include/material.hpp index d11bffb..46c9486 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" @@ -15,30 +15,29 @@ struct lambertian_material { lambertian_material(const texture_t& a) : albedo { a } {} - bool scatter(auto& ctx, const ray& r_in, const hit_record& rec, - color& attenuation, ray& scattered) const { - auto& rng = ctx.rng; + 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&& arg) { return arg.value(ctx, rec); }, albedo); + attenuation *= dev_visit([&](auto&& t) { return t.value(ctx, rec); }, 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; }; 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) } {} - bool scatter(auto& ctx, const ray& r_in, const hit_record& rec, - color& attenuation, ray& scattered) const { - auto& rng = ctx.rng; + 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); 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 @@ -49,7 +48,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 { @@ -65,20 +64,19 @@ struct dielectric_material { 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 { + 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 - auto& rng = ctx.rng; + LocalPseudoRNG rng { toseed(r_in.direction(), r_in.origin()) }; 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()) + 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([&](auto&& t) { return t.value(ctx, rec); }, emit); } texture_t emit; @@ -116,12 +114,11 @@ struct isotropic_material { isotropic_material(texture_t& a) : albedo { a } {} - bool scatter(auto& ctx, const ray& r_in, const hit_record& rec, - color& attenuation, ray& scattered) const { - auto& rng = ctx.rng; + 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&& arg) { return arg.value(ctx, rec); }, albedo); + attenuation *= dev_visit([&](auto&& t) { return t.value(ctx, rec); }, albedo); return true; } @@ -133,5 +130,4 @@ struct isotropic_material { using material_t = std::variant; - #endif 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/rectangle.hpp b/include/rectangle.hpp index 7ac9a08..f36cbd2 100644 --- a/include/rectangle.hpp +++ b/include/rectangle.hpp @@ -28,7 +28,7 @@ class xy_rect { , material_type { mat_type } {} /// Compute ray interaction with rectangle - 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; @@ -66,7 +66,7 @@ class xz_rect { , material_type { mat_type } {} /// Compute ray interaction with rectangle - 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; @@ -104,7 +104,7 @@ class yz_rect { , material_type { mat_type } {} /// Compute ray interaction with rectangle - 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/render.hpp b/include/render.hpp index 565f037..91efdb6 100644 --- a/include/render.hpp +++ b/include/render.hpp @@ -3,29 +3,20 @@ #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, +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& rng = ctx.rng; auto get_color = [&](const ray& r) { auto hit_world = [&](const ray& r, hit_record& rec, material_t& material_type) { @@ -36,9 +27,9 @@ inline auto render_pixel(auto& ctx, int x_coord, int y_coord, camera const& cam, // 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); + [&](auto&& object) { + return object.hit(r, 0.001f, closest_so_far, temp_rec, + temp_material_type); }, hittable_acc[i])) { hit_anything = true; @@ -58,12 +49,11 @@ 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); }, + emitted = dev_visit([&](auto&& mat) { return mat.emitted(ctx, rec); }, material_type); if (dev_visit( - [&](auto&& arg) { - return arg.scatter(ctx, cur_ray, rec, cur_attenuation, - scattered); + [&](auto&& mat) { + return mat.scatter(ctx, cur_ray, rec, cur_attenuation, scattered); }, material_type)) { // On hitting the object, the ray gets scattered @@ -90,12 +80,16 @@ 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 }; }; - + 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.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); } @@ -107,17 +101,16 @@ 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, +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& texture_acc) { if constexpr (buildparams::use_single_task) { cgh.single_task([=] { - LocalPseudoRNG rng; - task_context ctx { rng, texture_acc.get_pointer() }; + 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( - ctx, 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 { @@ -127,21 +120,17 @@ 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, texture_acc.get_pointer() }; - render_pixel( - ctx, x_coord, y_coord, cam_ptr, hittable_acc, fb_acc); + task_context ctx { texture_acc.get_pointer() }; + render_pixel(ctx, 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)); @@ -153,8 +142,6 @@ void render(sycl::queue& queue, sycl::buffer& frame_buf, 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(width, height, depth, samples, cgh, cam, hittables_acc, fb_acc, texture_acc); }); } diff --git a/include/rtweekend.hpp b/include/rtweekend.hpp index db7618d..af42e2c 100644 --- a/include/rtweekend.hpp +++ b/include/rtweekend.hpp @@ -3,7 +3,9 @@ #include #include +#include #include +#include #include #include #include @@ -16,19 +18,49 @@ #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; +} + +/** + @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)); + 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 & 31) << 21; + uint32_t shifted3 = (y1 & 63) << 15; + 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 { public: @@ -36,42 +68,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 +113,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 }; } @@ -97,7 +129,6 @@ class LocalPseudoRNG { kernel callees */ struct task_context { - LocalPseudoRNG rng; // See image_texture in texture.hpp for more details sycl::global_ptr texture_data; }; diff --git a/include/sphere.hpp b/include/sphere.hpp index ccc5b26..804bb02 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 @@ -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 50002cc..e0a42b0 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 @@ -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; } @@ -53,16 +53,12 @@ struct checker_texture { /** @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: @@ -88,9 +84,7 @@ struct image_texture { 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 */ @@ -118,9 +112,7 @@ struct image_texture { /** @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() { 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/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/include/visit.hpp b/include/visit.hpp index 9aa8e22..b1417ee 100644 --- a/include/visit.hpp +++ b/include/visit.hpp @@ -45,6 +45,18 @@ 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. diff --git a/src/main.cpp b/src/main.cpp index df7d03a..e7b6bf2 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -1,10 +1,13 @@ #include "sycl.hpp" +#include "xorshift.hpp" #include #include #include +#include #include #include #include +#include #include #include @@ -30,7 +33,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(); @@ -40,7 +43,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( @@ -58,10 +60,26 @@ 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 || argc > 7) { + std::cerr << "Usage: sycl-rt OUT_WIDTH OUT_HEIGHT DEPTH SAMPLES " + "[SPHERE_INC [RAND_SEED]]" + << 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] }); + 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; @@ -73,14 +91,14 @@ int main() { 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 { 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) { // 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 +108,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 { @@ -129,11 +147,13 @@ int main() { 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 + // 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 }))); @@ -142,9 +162,8 @@ int main() { 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); - - // // Add a sphere with a SYCL logo in the background hittables.emplace_back( sphere { point { -60, 3, 5 }, 4, lambertian_material { t } }); @@ -182,13 +201,10 @@ int main() { aperture, focus_dist, 0.0f, 1.0f }; - // Sample per pixel - constexpr auto samples = 100; - // 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);