Diagnose unsupported overlaps between host- and command group accessors

include/simsycl/detail/parallel_for.hh

@@ -2,7 +2,6 @@
 
 #include "allocation.hh"
 
-#include "../schedule.hh"
 #include "../sycl/device.hh"
 #include "../sycl/forward.hh"
 #include "../sycl/id.hh"

include/simsycl/sycl/buffer.hh

@@ -7,7 +7,6 @@
 #include "../detail/allocation.hh"
 #include "../detail/reference_type.hh"
 
-#include <concepts>
 #include <cstring>
 #include <memory>
 #include <mutex>
@@ -76,8 +75,60 @@ constexpr bool is_container_v = requires(C &c) {
 template<typename C, typename T>
 concept Container = is_container_v<C, T>;
 
+template<int Dimensions>
+struct accessed_range {
+    sycl::id<Dimensions> offset;
+    sycl::range<Dimensions> range;
+    sycl::access_mode mode;
+
+    accessed_range(sycl::id<Dimensions> offset, sycl::range<Dimensions> range, sycl::access_mode mode)
+        : offset(offset), range(range), mode(mode) {}
+
+    friend bool operator==(const accessed_range &lhs, const accessed_range &rhs) = default;
+
+    bool conflicts_with(const accessed_range &other) const {
+        if(mode == sycl::access_mode::read && other.mode == sycl::access_mode::read) return false;
+
+        for(int i = 0; i < Dimensions; ++i) {
+            if(offset[i] < other.offset[i]) {
+                if(offset[i] + range[i] <= other.offset[i]) { return false; }
+            } else {
+                if(other.offset[i] + other.range[i] <= offset[i]) { return false; }
+            }
+        }
+        return true;
+    }
+};
+
+// Base class for buffer_state necessary to keep a reference in accessor instances which do not know AllocatorT
+template<int Dimensions>
+struct buffer_access_validator {
+    std::vector<accessed_range<Dimensions>> live_host_accesses;
+
+    buffer_access_validator() = default;
+    buffer_access_validator(const buffer_access_validator &) = delete;
+    buffer_access_validator(buffer_access_validator &&) = delete;
+    buffer_access_validator &operator=(const buffer_access_validator &) = delete;
+    buffer_access_validator &operator=(buffer_access_validator &&) = delete;
+
+    void begin_host_access(const detail::accessed_range<Dimensions> &range) { live_host_accesses.push_back(range); }
+
+    void end_host_access(const detail::accessed_range<Dimensions> &range) {
+        auto &live = live_host_accesses;
+        live.erase(std::remove(live.begin(), live.end(), range), live.end());
+    }
+
+    void check_access_from_command_group(const detail::accessed_range<Dimensions> &range) const {
+        for(const auto &live_range : live_host_accesses) {
+            SIMSYCL_CHECK(!live_range.conflicts_with(range)
+                && "Command group accessor overlaps with a live host accessor for the same buffer range, this is not "
+                   "supported by SimSYCL unless both are read-only accesses");
+        }
+    }
+};
+
 template<typename T, int Dimensions, typename AllocatorT>
-struct buffer_state {
+struct buffer_state : buffer_access_validator<Dimensions> {
     using write_back_fn = std::function<void(const T *, size_t)>;
 
     sycl::range<Dimensions> range;
@@ -291,6 +342,9 @@ class buffer final : public detail::reference_type<buffer<T, Dimensions, Allocat
     template<typename U, int D, typename A>
     friend U *simsycl::detail::get_buffer_data(sycl::buffer<U, D, A> &buf);
 
+    template<typename U, int D, typename A>
+    friend detail::buffer_access_validator<D> &detail::get_buffer_access_validator(const sycl::buffer<U, D, A> &buf);
+
     using reference_type::state;
 
     static write_back_fn write_back_to(T out) {
@@ -357,4 +411,9 @@ T *get_buffer_data(sycl::buffer<T, Dimensions, AllocatorT> &buf) {
     return buf.state().buffer;
 }
 
+template<typename T, int Dimensions, typename AllocatorT>
+buffer_access_validator<Dimensions> &get_buffer_access_validator(const sycl::buffer<T, Dimensions, AllocatorT> &buf) {
+    return buf.state();
+}
+
 } // namespace simsycl::detail

include/simsycl/sycl/forward.hh

@@ -162,9 +162,15 @@ concurrent_sub_group &get_concurrent_group(const sycl::sub_group &g);
 template<int Dimensions>
 concurrent_group &get_concurrent_group(const sycl::group<Dimensions> &g);
 
+template<int Dimensions>
+struct buffer_access_validator;
+
 template<typename T, int Dimensions, typename AllocatorT>
 T *get_buffer_data(sycl::buffer<T, Dimensions, AllocatorT> &buf);
 
+template<typename T, int Dimensions, typename AllocatorT>
+buffer_access_validator<Dimensions> &get_buffer_access_validator(const sycl::buffer<T, Dimensions, AllocatorT> &buf);
+
 sycl::handler make_handler(const sycl::device &device);
 
 sycl::interop_handle make_interop_handle();

src/simsycl/device.cc

@@ -16,9 +16,7 @@ struct device_state {
     weak_ref<sycl::device> parent;
 };
 
-size_t *device_bytes_free(const sycl::device &device) {
-    return &device.state().bytes_free;
-}
+size_t *device_bytes_free(const sycl::device &device) { return &device.state().bytes_free; }
 
 int default_selector::operator()(const sycl::device &device) const {
     return device.is_gpu() || device.is_accelerator() ? 1 : 0;
@@ -473,7 +471,7 @@ SIMSYCL_DETAIL_DEPRECATED_IN_SYCL bool device::has_extension(const std::string &
 
 std::vector<device> device::get_devices(info::device_type type) {
     auto &devices = detail::get_devices();
-    if (type == info::device_type::all) return devices;
+    if(type == info::device_type::all) return devices;
 
     std::vector<device> result;
     std::copy_if(devices.begin(), devices.end(), std::back_inserter(result),

test/check_tests.cc

@@ -1,8 +1,11 @@
-#include <catch2/catch_test_macros.hpp>
-#include <catch2/matchers/catch_matchers_string.hpp>
+#include "test_utils.hh"
 
 #include <sycl/sycl.hpp>
 
+#include <catch2/catch_template_test_macros.hpp>
+#include <catch2/catch_test_macros.hpp>
+#include <catch2/matchers/catch_matchers_string.hpp>
+
 using namespace simsycl;
 using Catch::Matchers::ContainsSubstring;
 
@@ -26,9 +29,11 @@ TEST_CASE("SIMSYCL_CHECK follows the configured setting - LOG", "[check]") {
     std::cout.rdbuf(stdout_buffer);
     REQUIRE_THAT(oss.str(), ContainsSubstring("SimSYCL check failed: false && \"Bla\" at "));
 }
+
 #endif
 
 #if SIMSYCL_CHECK_MODE == SIMSYCL_CHECK_THROW
+
 TEST_CASE("SIMSYCL_CHECK follows the configured setting - THROW", "[check]") {
     REQUIRE_THROWS_WITH(
         [] { SIMSYCL_CHECK(false && "Bla"); }(), ContainsSubstring("SimSYCL check failed: false && \"Bla\" at "));
@@ -41,4 +46,58 @@ TEST_CASE("Exceptions are propagated out of work items", "[check][exceptions]")
     }),
         ContainsSubstring("SimSYCL check failed: false && \"Bla\" at "));
 }
-#endif
+
+template<typename T>
+extern const sycl::access_mode accessor_mode_v;
+
+template<typename T, int D, sycl::access_mode M, sycl::target A>
+constexpr sycl::access_mode accessor_mode_v<sycl::accessor<T, D, M, A>> = M;
+
+template<typename T, int D, sycl::access_mode M>
+constexpr sycl::access_mode accessor_mode_v<sycl::host_accessor<T, D, M>> = M;
+
+using host_accessor_types = std::tuple<                                               //
+    sycl::host_accessor<int, 1, sycl::access_mode::read>,                             //
+    sycl::host_accessor<int, 0, sycl::access_mode::read>,                             //
+    sycl::host_accessor<int, 1, sycl::access_mode::read_write>,                       //
+    sycl::host_accessor<int, 0, sycl::access_mode::read_write>,                       //
+    sycl::accessor<int, 1, sycl::access_mode::read, sycl::target::host_buffer>,       //
+    sycl::accessor<int, 0, sycl::access_mode::read, sycl::target::host_buffer>,       //
+    sycl::accessor<int, 1, sycl::access_mode::read_write, sycl::target::host_buffer>, //
+    sycl::accessor<int, 0, sycl::access_mode::read_write, sycl::target::host_buffer>>;
+
+using command_group_accessor_types = std::tuple<                                        //
+    sycl::accessor<int, 1, sycl::access_mode::read, sycl::target::global_buffer>,       //
+    sycl::accessor<int, 0, sycl::access_mode::read, sycl::target::global_buffer>,       //
+    sycl::accessor<int, 1, sycl::access_mode::read_write, sycl::target::global_buffer>, //
+    sycl::accessor<int, 0, sycl::access_mode::read_write, sycl::target::global_buffer>, //
+    sycl::accessor<int, 1, sycl::access_mode::read, sycl::target::constant_buffer>,     //
+    sycl::accessor<int, 0, sycl::access_mode::read, sycl::target::constant_buffer>>;
+
+using accessor_type_combinations = test::tuple_cross_product<host_accessor_types, command_group_accessor_types>::type;
+
+TEMPLATE_LIST_TEST_CASE("Overlapping lifetimes between host- and command-group accessors are diagnosed", "[check]",
+    accessor_type_combinations) {
+    using host_accessor_type = std::tuple_element_t<0, TestType>;
+    using command_group_accessor_type = std::tuple_element_t<1, TestType>;
+
+    sycl::buffer<int, 1> buf(100);
+    host_accessor_type host_acc(buf);
+
+    const auto submit_overlapping_command_group = [&] {
+        sycl::queue().submit([&](sycl::handler &cgh) {
+            command_group_accessor_type acc(buf, cgh);
+            cgh.single_task([=] { (void)acc; });
+        });
+    };
+
+    if(accessor_mode_v<host_accessor_type> == sycl::access_mode::read
+        && accessor_mode_v<command_group_accessor_type> == sycl::access_mode::read) {
+        REQUIRE_NOTHROW(submit_overlapping_command_group());
+    } else {
+        REQUIRE_THROWS_WITH(
+            submit_overlapping_command_group(), ContainsSubstring("overlaps with a live host accessor"));
+    }
+}
+
+#endif

test/group_op_tests.cc

@@ -1,6 +1,7 @@
 #include <catch2/catch_test_macros.hpp>
 #include <catch2/generators/catch_generators.hpp>
 
+#include <simsycl/schedule.hh>
 #include <sycl/sycl.hpp>
 
 #include "test_utils.hh"
@@ -14,9 +15,9 @@ void check_group_op_sequence(const G &g, const std::vector<detail::group_operati
     for(size_t i = 0; i < expected_ids.size(); ++i) { CHECK(group_instance.operations[i].id == expected_ids[i]); }
 }
 
-#define REPEAT_FOR_ALL_SCHEDULES \
-    std::string schedule = GENERATE(values<std::string>({"round_robin", "shuffle"})); \
-    CAPTURE(schedule); \
+#define REPEAT_FOR_ALL_SCHEDULES                                                                                       \
+    std::string schedule = GENERATE(values<std::string>({"round_robin", "shuffle"}));                                  \
+    CAPTURE(schedule);                                                                                                 \
     if(schedule == "shuffle") { set_cooperative_schedule(std::make_unique<shuffle_schedule>()); }
 
 TEST_CASE("Group barriers behave as expected", "[group_op]") {

test/simulation_tests.cc

@@ -1,4 +1,5 @@
 #include "test_utils.hh"
+#include <simsycl/schedule.hh>
 #include <sycl/sycl.hpp>
 
 #include <catch2/catch_test_macros.hpp>

test/test_utils.hh

@@ -3,6 +3,8 @@
 #include <simsycl/sycl/vec.hh>
 #include <simsycl/system.hh>
 
+#include <tuple>
+
 namespace simsycl::test {
 
 template<typename DeviceSetup>
@@ -20,4 +22,19 @@ inline bool check_bool_vec(simsycl::sycl::vec<bool, Dimensions> a) {
     return true;
 }
 
+// From https://stackoverflow.com/a/70405002/1522056
+template<typename T1, typename T2>
+class tuple_cross_product {
+    template<typename T, typename... Ts>
+    static auto inner_helper(T &&, std::tuple<Ts...> &&)
+        -> decltype(std::make_tuple(std::make_tuple(std::declval<T>(), std::declval<Ts>())...));
+
+    template<typename... Ts, typename T>
+    static auto outer_helper(std::tuple<Ts...> &&, T &&)
+        -> decltype(std::tuple_cat(inner_helper(std::declval<Ts>(), std::declval<T>())...));
+
+  public:
+    using type = decltype(outer_helper(std::declval<T1>(), std::declval<T2>()));
+};
+
 }; // namespace simsycl::test