diff --git a/include/alpaka/dev/DevGenericSycl.hpp b/include/alpaka/dev/DevGenericSycl.hpp index 85e00738cf98..be783b658394 100644 --- a/include/alpaka/dev/DevGenericSycl.hpp +++ b/include/alpaka/dev/DevGenericSycl.hpp @@ -70,7 +70,7 @@ namespace alpaka for(auto& q_ptr : m_queues) { if(auto ptr = q_ptr.lock(); ptr != nullptr) - ptr->register_dependency(event); + ptr->getNativeHandle().ext_oneapi_submit_barrier({event}); } } diff --git a/include/alpaka/event/EventGenericSycl.hpp b/include/alpaka/event/EventGenericSycl.hpp index 68011a0247cd..800bb20fc8d9 100644 --- a/include/alpaka/event/EventGenericSycl.hpp +++ b/include/alpaka/event/EventGenericSycl.hpp @@ -87,7 +87,7 @@ namespace alpaka::trait { static auto enqueue(QueueGenericSyclNonBlocking& queue, EventGenericSycl& event) { - event.setEvent(queue.m_spQueueImpl->get_last_event()); + event.setEvent(queue.getNativeHandle().ext_oneapi_submit_barrier()); } }; @@ -97,7 +97,7 @@ namespace alpaka::trait { static auto enqueue(QueueGenericSyclBlocking& queue, EventGenericSycl& event) { - event.setEvent(queue.m_spQueueImpl->get_last_event()); + event.setEvent(queue.getNativeHandle().ext_oneapi_submit_barrier()); } }; @@ -120,7 +120,7 @@ namespace alpaka::trait { static auto waiterWaitFor(QueueGenericSyclNonBlocking& queue, EventGenericSycl const& event) { - queue.m_spQueueImpl->register_dependency(event.getNativeHandle()); + queue.getNativeHandle().ext_oneapi_submit_barrier({event.getNativeHandle()}); } }; @@ -130,7 +130,7 @@ namespace alpaka::trait { static auto waiterWaitFor(QueueGenericSyclBlocking& queue, EventGenericSycl const& event) { - queue.m_spQueueImpl->register_dependency(event.getNativeHandle()); + queue.getNativeHandle().ext_oneapi_submit_barrier({event.getNativeHandle()}); } }; diff --git a/include/alpaka/mem/buf/sycl/Copy.hpp b/include/alpaka/mem/buf/sycl/Copy.hpp index 2ba963935851..562371bdaacf 100644 --- a/include/alpaka/mem/buf/sycl/Copy.hpp +++ b/include/alpaka/mem/buf/sycl/Copy.hpp @@ -93,7 +93,7 @@ namespace alpaka::detail using TaskCopySyclBase::TaskCopySyclBase; - auto operator()(sycl::queue& queue, std::vector const& requirements) const -> sycl::event + auto operator()(sycl::queue& queue) const -> sycl::event { ALPAKA_DEBUG_MINIMAL_LOG_SCOPE; @@ -126,8 +126,7 @@ namespace alpaka::detail this->m_srcMemNative + (castVec(idx) * srcPitchBytesWithoutOutmost) .foldrAll(std::plus())), - static_cast(this->m_extentWidthBytes), - requirements)); + static_cast(this->m_extentWidthBytes))); }); } @@ -144,7 +143,7 @@ namespace alpaka::detail using TaskCopySyclBase, TViewDst, TViewSrc, TExtent>::TaskCopySyclBase; using Elem = alpaka::Elem; - auto operator()(sycl::queue& queue, std::vector const& requirements) const -> sycl::event + auto operator()(sycl::queue& queue) const -> sycl::event { ALPAKA_DEBUG_MINIMAL_LOG_SCOPE; @@ -156,8 +155,7 @@ namespace alpaka::detail return queue.memcpy( reinterpret_cast(this->m_dstMemNative), reinterpret_cast(this->m_srcMemNative), - sizeof(Elem) * static_cast(this->m_extent.prod()), - requirements); + sizeof(Elem) * static_cast(this->m_extent.prod())); } else { @@ -187,9 +185,9 @@ namespace alpaka::detail ALPAKA_ASSERT(getExtentVec(viewSrc).prod() == 1u); } - auto operator()(sycl::queue& queue, std::vector const& requirements) const -> sycl::event + auto operator()(sycl::queue& queue) const -> sycl::event { - return queue.memcpy(m_dstMemNative, m_srcMemNative, sizeof(Elem), requirements); + return queue.memcpy(m_dstMemNative, m_srcMemNative, sizeof(Elem)); } void* m_dstMemNative; diff --git a/include/alpaka/mem/buf/sycl/Set.hpp b/include/alpaka/mem/buf/sycl/Set.hpp index 58712cfd167c..aa8c9bee3842 100644 --- a/include/alpaka/mem/buf/sycl/Set.hpp +++ b/include/alpaka/mem/buf/sycl/Set.hpp @@ -83,7 +83,7 @@ namespace alpaka using TaskSetSyclBase::TaskSetSyclBase; - auto operator()(sycl::queue& queue, std::vector const& requirements) const -> sycl::event + auto operator()(sycl::queue& queue) const -> sycl::event { ALPAKA_DEBUG_MINIMAL_LOG_SCOPE; @@ -112,8 +112,7 @@ namespace alpaka + (castVec(idx) * dstPitchBytesWithoutOutmost) .foldrAll(std::plus())), this->m_byte, - static_cast(this->m_extentWidthBytes), - requirements)); + static_cast(this->m_extentWidthBytes))); }); } @@ -128,7 +127,7 @@ namespace alpaka { using TaskSetSyclBase, TView, TExtent>::TaskSetSyclBase; - auto operator()(sycl::queue& queue, std::vector const& requirements) const -> sycl::event + auto operator()(sycl::queue& queue) const -> sycl::event { ALPAKA_DEBUG_MINIMAL_LOG_SCOPE; @@ -140,8 +139,7 @@ namespace alpaka return queue.memset( reinterpret_cast(this->m_dstMemNative), this->m_byte, - static_cast(this->m_extentWidthBytes), - requirements); + static_cast(this->m_extentWidthBytes)); } else { @@ -178,14 +176,14 @@ namespace alpaka } # endif - auto operator()(sycl::queue& queue, std::vector const& requirements) const -> sycl::event + auto operator()(sycl::queue& queue) const -> sycl::event { ALPAKA_DEBUG_MINIMAL_LOG_SCOPE; # if ALPAKA_DEBUG >= ALPAKA_DEBUG_FULL printDebug(); # endif - return queue.memset(reinterpret_cast(m_dstMemNative), m_byte, sizeof(Elem), requirements); + return queue.memset(reinterpret_cast(m_dstMemNative), m_byte, sizeof(Elem)); } std::uint8_t const m_byte; diff --git a/include/alpaka/queue/sycl/QueueGenericSyclBase.hpp b/include/alpaka/queue/sycl/QueueGenericSyclBase.hpp index ca62f7250831..d60cb62a5aef 100644 --- a/include/alpaka/queue/sycl/QueueGenericSyclBase.hpp +++ b/include/alpaka/queue/sycl/QueueGenericSyclBase.hpp @@ -13,8 +13,6 @@ #include #include #include -#include -#include #include #include #include @@ -73,36 +71,9 @@ namespace alpaka::detail } } - // Don't call this without locking first! - auto clean_dependencies() -> void - { - // Clean up completed events - auto const start = std::begin(m_dependencies); - auto const old_end = std::end(m_dependencies); - auto const new_end = std::remove_if( - start, - old_end, - [](sycl::event ev) { - return ev.get_info() - == sycl::info::event_command_status::complete; - }); - - m_dependencies.erase(new_end, old_end); - } - - auto register_dependency(sycl::event event) -> void - { - std::lock_guard lock{m_mutex}; - - clean_dependencies(); - m_dependencies.push_back(event); - } - auto empty() const -> bool { - std::shared_lock lock{m_mutex}; - return m_last_event.get_info() - == sycl::info::event_command_status::complete; + return m_queue.ext_oneapi_empty(); } auto wait() -> void @@ -111,41 +82,26 @@ namespace alpaka::detail m_queue.wait_and_throw(); } - auto get_last_event() const -> sycl::event - { - std::shared_lock lock{m_mutex}; - return m_last_event; - } - template auto enqueue(TTask const& task) -> void { { - std::lock_guard lock{m_mutex}; - - clean_dependencies(); - // Execute task if constexpr(is_sycl_task && !is_sycl_kernel) // Copy / Fill { - m_last_event = task(m_queue, m_dependencies); // Will call queue.{copy, fill} internally + task(m_queue); // Will call queue.{copy, fill} internally } else { - m_last_event = m_queue.submit( + m_queue.submit( [this, &task](sycl::handler& cgh) { - if(!m_dependencies.empty()) - cgh.depends_on(m_dependencies); - if constexpr(is_sycl_kernel) // Kernel task(cgh); // Will call cgh.parallel_for internally else // Host cgh.host_task(task); }); } - - m_dependencies.clear(); } if constexpr(TBlocking) @@ -157,10 +113,6 @@ namespace alpaka::detail return m_queue; } - std::vector m_dependencies; - sycl::event m_last_event; - std::shared_mutex mutable m_mutex; - private: sycl::queue m_queue; };