diff --git a/include/camp/resource/sycl.hpp b/include/camp/resource/sycl.hpp index 4206f29..90dc074 100644 --- a/include/camp/resource/sycl.hpp +++ b/include/camp/resource/sycl.hpp @@ -34,19 +34,35 @@ namespace resources class SyclEvent { public: - // TODO: make this actually work - SyclEvent(sycl::queue& CAMP_UNUSED_ARG(qu)) { m_event = sycl::event(); } + explicit SyclEvent(sycl::event e) + : m_event(std::move(e)) + {} + + // TODO: see what overhead an empty submit has + SyclEvent(sycl::queue& qu) + { + if (!qu.is_in_order()) { + ::camp::throw_re("Queue is not in_order."); + } + m_event = qu.submit([&](::sycl::handler& CAMP_UNUSED_ARG(h)) {}); + } SyclEvent(Sycl& res); - bool check() const { return true; } + bool check() const + { + return m_event.get_info() + == sycl::info::event_command_status::complete; + } + + void wait() const { m_event.wait(); } // sycl::event::wait is non-const - void wait() const { getSyclEvent_t().wait(); } + sycl::event& getSyclEvent_t() { return m_event; } - sycl::event getSyclEvent_t() const { return m_event; } + sycl::event const& getSyclEvent_t() const { return m_event; } private: - sycl::event m_event; + mutable sycl::event m_event; // mutable as use non-const member function }; class Sycl @@ -265,7 +281,9 @@ namespace resources { auto* sycl_event = e->try_get(); if (sycl_event) { - (sycl_event->getSyclEvent_t()).wait(); + qu.submit([&](::sycl::handler& h) { + h.depends_on(sycl_event->getSyclEvent_t()); + }); } else { e->wait(); } @@ -356,7 +374,7 @@ namespace resources inline SyclEvent::SyclEvent(Sycl &res) : SyclEvent(res.get_queue()) - { } + {} } // namespace v1 diff --git a/test/resource.cpp b/test/resource.cpp index 55241a0..de134d6 100644 --- a/test/resource.cpp +++ b/test/resource.cpp @@ -658,6 +658,63 @@ TEST(CampResource, Wait) #endif } +template +void test_event_wait() +{ + auto r = Res(); + const auto typed_event = r.get_event(); + const Event event = r.get_event_erased(); + typed_event.wait(); + event.wait(); +} + +// +TEST(CampEvent, Wait) +{ + test_event_wait(); +#ifdef CAMP_HAVE_CUDA + test_event_wait(); +#endif +#ifdef CAMP_HAVE_HIP + test_event_wait(); +#endif +#ifdef CAMP_HAVE_OMP_OFFLOAD + test_event_wait(); +#endif +#ifdef CAMP_HAVE_SYCL + test_event_wait(); +#endif +} + +template +void test_event_check() +{ + auto r = Res(); + const auto typed_event = r.get_event(); + const Event event = r.get_event_erased(); + // checking in a loop should always eventually return true + while (!typed_event.check()) {} + while (!event.check()) {} +} + +// +TEST(CampEvent, Check) +{ + test_event_check(); +#ifdef CAMP_HAVE_CUDA + test_event_check(); +#endif +#ifdef CAMP_HAVE_HIP + test_event_check(); +#endif +#ifdef CAMP_HAVE_OMP_OFFLOAD + test_event_check(); +#endif +#ifdef CAMP_HAVE_SYCL + test_event_check(); +#endif +} + template void test_concrete_resource_trait() {