Alternative work distribution

CMakeLists.txt

@@ -1,6 +1,7 @@
 cmake_minimum_required(VERSION 3.16)
 
-option(USE_SINGLE_TASK "Use a SYCL executor based on .single_task() instead of .parallel_for(), better for FPGA" OFF)
+option(USE_SINGLE_TASK "Use a SYCL executor that loops over pixel in one task instead of using a parallel_for(), better for FPGA" OFF)
+option(SANITIZE_THREADS "Activate thread sanitizer" OFF)
 set(SYCL_CXX_COMPILER "" CACHE STRING "Path to the SYCL compiler. Defaults to using triSYCL CPU implementation" )
 # Use SYCL host device by default
 set(SYCL_DEVICE_TRIPLE "" CACHE STRING "Device triple to be used. only used with SYCL_CXX_COMPILER")
@@ -10,27 +11,30 @@ if (NOT "${SYCL_CXX_COMPILER}" STREQUAL "")
 endif()
 
 project(SYCL-path-tracer LANGUAGES CXX)
+set(CMAKE_CXX_STANDARD 20)
+set(CMAKE_CXX_STANDARD_REQUIRED ON)
+set(CMAKE_CXX_EXTENSIONS OFF)
 
 # Use triSYCL
 if ("${SYCL_CXX_COMPILER}" STREQUAL "")
   include(FindtriSYCL)
 else()
   set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -fsycl -Xsycl-target-frontend -fno-exceptions")
   if (NOT "${SYCL_DEVICE_TRIPLE}" STREQUAL "")
-    message(STATUS "targeting: ${SYCL_DEVICE_TRIPLE}")
-    set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -fsycl-targets=${SYCL_DEVICE_TRIPLE}")
+	message(STATUS "targeting: ${SYCL_DEVICE_TRIPLE}")
+	set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -fsycl-targets=${SYCL_DEVICE_TRIPLE}")
   endif()
   set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DUSE_SYCL_COMPILER")
 endif()
 
- 
+
 # Set a default build type if none was specified
 if(NOT CMAKE_BUILD_TYPE AND NOT CMAKE_CONFIGURATION_TYPES)
   message(STATUS "Setting build type to Release as none was specified.")
   set(CMAKE_BUILD_TYPE "Release" CACHE
-      STRING "Choose the type of build." FORCE)
+	  STRING "Choose the type of build." FORCE)
   set_property(CACHE CMAKE_BUILD_TYPE PROPERTY STRINGS
-    "Debug" "Release" "MinSizeRel" "RelWithDebInfo")
+	"Debug" "Release" "MinSizeRel" "RelWithDebInfo")
 endif()
 
 
@@ -49,14 +53,16 @@ endif()
 # Use C+20
 target_compile_features(sycl-rt PRIVATE cxx_std_20)
 
+if (SANITIZE_THREADS)
+target_compile_options(sycl-rt PRIVATE
+                       -fno-omit-frame-pointer -fsanitize=thread)
+target_link_options(sycl-rt PRIVATE -fsanitize=thread)
+endif()
 # To use various code sanitizer:
 #target_compile_options(sycl-rt PRIVATE
 #                       -fno-omit-frame-pointer -fsanitize=address)
 #target_link_options(sycl-rt PRIVATE -fsanitize=address)
 #target_compile_options(sycl-rt PRIVATE
-#                       -fno-omit-frame-pointer -fsanitize=thread)
-#target_link_options(sycl-rt PRIVATE -fsanitize=thread)
-#target_compile_options(sycl-rt PRIVATE
 #                       -fno-omit-frame-pointer -fsanitize=undefined)
 #target_link_options(sycl-rt PRIVATE -fsanitize=undefined)
 #target_compile_options(sycl-rt PRIVATE
@@ -67,8 +73,9 @@ target_compile_features(sycl-rt PRIVATE cxx_std_20)
 if(USE_SINGLE_TASK)
   # On FPGA use a loop on image pixels instead of a parallel_for
   set_property(TARGET sycl-rt
-    APPEND PROPERTY
-    COMPILE_DEFINITIONS USE_SINGLE_TASK=)
+	APPEND PROPERTY
+	COMPILE_DEFINITIONS USE_SINGLE_TASK=)
 endif()
 
 message(STATUS "path_tracer USE_SINGLE_TASK:      ${USE_SINGLE_TASK}")
+message(STATUS "path_tracer SANITIZE_THREADS:      ${SANITIZE_THREADS}")

include/box.hpp

@@ -13,40 +13,40 @@ class box {
   /// 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);
+	  : 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;
+		   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;

include/build_parameters.hpp

@@ -1,11 +1,18 @@
 #ifndef BUILD_PARAMETERS_HPP
 #define BUILD_PARAMETERS_HPP
 namespace buildparams {
+
 #if USE_SINGLE_TASK
 constexpr bool use_single_task = true;
 #else
 constexpr bool use_single_task = false;
 #endif
+
+#if USE_SYCL_COMPILER
+constexpr bool use_sycl_compiler = true;
+#else
+constexpr bool use_sycl_compiler = false;
+#endif
 }
 
 

include/camera.hpp

@@ -8,12 +8,12 @@
 
 /** Camera model
 
-    This implements:
+	This implements:
 
-    -
+	-
    https://raytracing.github.io/books/RayTracingInOneWeekend.html#positionablecamera
 
-    - https://raytracing.github.io/books/RayTracingInOneWeekend.html#defocusblur
+	- https://raytracing.github.io/books/RayTracingInOneWeekend.html#defocusblur
 */
 class camera {
 
@@ -47,55 +47,55 @@ class camera {
  public:
   /** Create a parameterized camera
 
-      \param[in] look_from is the position of the camera
+	  \param[in] look_from is the position of the camera
 
-      \param[in] look_at is a point the camera is looking at
+	  \param[in] look_at is a point the camera is looking at
 
-      \param[in] vup is the “view up” orientation for the
-      camera. {0,1,0} means the usual vertical orientation
+	  \param[in] vup is the “view up” orientation for the
+	  camera. {0,1,0} means the usual vertical orientation
 
-      \param[in] degree_vfov is the vertical field-of-view in degrees
+	  \param[in] degree_vfov is the vertical field-of-view in degrees
 
-      \param[in] aspect_ratio is the ratio between the camera image
-      width and the camera image height
+	  \param[in] aspect_ratio is the ratio between the camera image
+	  width and the camera image height
 
-      \param[in] aperture is the lens aperture of the camera
+	  \param[in] aperture is the lens aperture of the camera
 
-      \param[in] focus_dist is the focus distance
+	  \param[in] focus_dist is the focus distance
   */
   camera(const point& look_from, const point& look_at, const vec& vup,
-         real_t degree_vfov, real_t aspect_ratio, real_t aperture,
-         real_t focus_dist, real_t _time0 = 0, real_t _time1 = 0)
-      : origin { look_from } {
-    auto theta = degrees_to_radians(degree_vfov);
-    auto h = std::tan(theta / 2);
-    auto viewport_height = 2.0f * h;
-    auto viewport_width = aspect_ratio * viewport_height;
-
-    w = unit_vector(look_from - look_at);
-    u = unit_vector(sycl::cross(vup, w));
-    v = sycl::cross(w, u);
-
-    horizontal = focus_dist * viewport_width * u;
-    vertical = focus_dist * viewport_height * v;
-    lower_left_corner = origin - horizontal / 2 - vertical / 2 - focus_dist * w;
-
-    lens_radius = aperture / 2;
-    time0 = _time0;
-    time1 = _time1;
+		 real_t degree_vfov, real_t aspect_ratio, real_t aperture,
+		 real_t focus_dist, real_t _time0 = 0, real_t _time1 = 0)
+	  : origin { look_from } {
+	auto theta = degrees_to_radians(degree_vfov);
+	auto h = std::tan(theta / 2);
+	auto viewport_height = 2.0f * h;
+	auto viewport_width = aspect_ratio * viewport_height;
+
+	w = unit_vector(look_from - look_at);
+	u = unit_vector(sycl::cross(vup, w));
+	v = sycl::cross(w, u);
+
+	horizontal = focus_dist * viewport_width * u;
+	vertical = focus_dist * viewport_height * v;
+	lower_left_corner = origin - horizontal / 2 - vertical / 2 - focus_dist * w;
+
+	lens_radius = aperture / 2;
+	time0 = _time0;
+	time1 = _time1;
   }
 
   /** Computes ray from camera passing through
-      viewport local coordinates (s,t) based on viewport
-      width, height and focus distance
+	  viewport local coordinates (s,t) based on viewport
+	  width, height and focus distance
   */
   ray get_ray(real_t s, real_t t) const {
-    vec rd = lens_radius * random_in_unit_disk();
-    vec offset = u * rd.x() + v * rd.y();
-    return { origin + offset,
-             lower_left_corner + s * horizontal + t * vertical - origin -
-                 offset,
-             random_float(time0, time1) };
+	vec rd = lens_radius * random_in_unit_disk();
+	vec offset = u * rd.x() + v * rd.y();
+	return { origin + offset,
+			 lower_left_corner + s * horizontal + t * vertical - origin -
+				 offset,
+			 random_float(time0, time1) };
   }
 };
 

include/constant_medium.hpp

@@ -16,63 +16,63 @@ using hittableVolume_t = std::variant<sphere, box>;
 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 } } {}
+	  : 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 } } {}
+	  : 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;
-    }
+		   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 (!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;
+	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 * log(random_float());
+	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 * log(random_float());
 
-    /// With lower density, hit_distance has higher probabilty
-    /// of being greater than distance_inside_boundary
-    if (hit_distance > distance_inside_boundary)
-      return false;
+	/// 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.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;
+	rec.normal = vec { 1, 0, 0 }; // arbitrary
+	rec.front_face = true;        // also arbitrary
+	return true;
   }
 
   hittableVolume_t boundary;