From c1af0b996c0bb841be72322b3ebad9d41f01b658 Mon Sep 17 00:00:00 2001 From: Hugh Delaney Date: Wed, 5 Oct 2022 14:38:11 +0100 Subject: [PATCH 1/5] WIP --- sycl/include/sycl/detail/cg_types.hpp | 4 ++ sycl/include/sycl/detail/property_helper.hpp | 4 +- sycl/include/sycl/handler.hpp | 7 +-- sycl/include/sycl/interop_handle.hpp | 29 +++++++++++- .../sycl/properties/all_properties.hpp | 1 + .../sycl/properties/host_task_properties.hpp | 46 +++++++++++++++++++ sycl/source/detail/scheduler/commands.cpp | 28 ++++++----- 7 files changed, 102 insertions(+), 17 deletions(-) create mode 100644 sycl/include/sycl/properties/host_task_properties.hpp diff --git a/sycl/include/sycl/detail/cg_types.hpp b/sycl/include/sycl/detail/cg_types.hpp index cef96878d2fe0..5bc6a9326feb6 100644 --- a/sycl/include/sycl/detail/cg_types.hpp +++ b/sycl/include/sycl/detail/cg_types.hpp @@ -228,6 +228,7 @@ class InteropTask { class HostTask { std::function MHostTask; std::function MInteropTask; + property_list MPropertyList; public: HostTask() : MHostTask([]() {}) {} @@ -238,6 +239,9 @@ class HostTask { void call() { MHostTask(); } void call(interop_handle handle) { MInteropTask(handle); } + + friend class sycl::handler; + friend class DispatchHostTask; }; // Class which stores specific lambda object. diff --git a/sycl/include/sycl/detail/property_helper.hpp b/sycl/include/sycl/detail/property_helper.hpp index ff0f4aa8568b0..0820b61e0d46a 100644 --- a/sycl/include/sycl/detail/property_helper.hpp +++ b/sycl/include/sycl/detail/property_helper.hpp @@ -35,8 +35,10 @@ enum DataLessPropKind { UseDefaultStream = 8, DiscardEvents = 9, DeviceReadOnly = 10, + HostTaskExecOnSubmit = 11, + HostTaskManualInteropSync = 12, // Indicates the last known dataless property. - LastKnownDataLessPropKind = 10, + LastKnownDataLessPropKind = 12, // Exceeding 32 may cause ABI breaking change on some of OSes. DataLessPropKindSize = 32 }; diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 335daae9904d6..9eec053224aea 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -1289,13 +1289,14 @@ class __SYCL_EXPORT handler { void()>::value || detail::check_fn_signature, void(interop_handle)>::value> - host_task_impl(FuncT &&Func) { + host_task_impl(FuncT &&Func, const property_list &PropList) { throwIfActionIsCreated(); MNDRDesc.set(range<1>(1)); MArgs = std::move(MAssociatedAccesors); MHostTask.reset(new detail::HostTask(std::move(Func))); + MHostTask->MPropertyList = std::move(PropList); setType(detail::CG::CodeplayHostTask); } @@ -1501,8 +1502,8 @@ class __SYCL_EXPORT handler { void()>::value || detail::check_fn_signature, void(interop_handle)>::value> - host_task(FuncT &&Func) { - host_task_impl(Func); + host_task(FuncT &&Func, const property_list &PropList = {}) { + host_task_impl(Func, PropList); } // replace _KERNELFUNCPARAM(KernelFunc) with KernelType KernelFunc diff --git a/sycl/include/sycl/interop_handle.hpp b/sycl/include/sycl/interop_handle.hpp index 8804073f827a9..5c1568a1f044c 100644 --- a/sycl/include/sycl/interop_handle.hpp +++ b/sycl/include/sycl/interop_handle.hpp @@ -14,6 +14,7 @@ #include #include #include +#include #include @@ -142,6 +143,25 @@ class interop_handle { #endif } + template + std::vector get_native_events() { +#ifndef __SYCL_DEVICE_ONLY__ + if (!MPropertyList + .has_property()) { + throw sycl::exception(make_error_code(errc::feature_not_supported), + "get_native_events can only be used in host task " + "with manual_interop_sync property"); + } + if (Backend != get_backend()) { + throw sycl::exception(make_error_code(errc::backend_mismatch), + "Incorrect backend argument was passed"); + } + return MRawEvents; +#endif + throw sycl::exception(make_error_code(errc::feature_not_supported), + "get_native_events should never be called on device"); + } + private: friend class detail::ExecCGCommand; friend class detail::DispatchHostTask; @@ -150,9 +170,12 @@ class interop_handle { interop_handle(std::vector MemObjs, const std::shared_ptr &Queue, const std::shared_ptr &Device, - const std::shared_ptr &Context) + const std::shared_ptr &Context, + std::vector RawEvents, + property_list PropList = {}) : MQueue(Queue), MDevice(Device), MContext(Context), - MMemObjs(std::move(MemObjs)) {} + MMemObjs(std::move(MemObjs)), MRawEvents(RawEvents), + MPropertyList(PropList) {} template backend_return_t> @@ -173,6 +196,8 @@ class interop_handle { std::shared_ptr MContext; std::vector MMemObjs; + std::vector MRawEvents; + property_list MPropertyList; }; } // __SYCL_INLINE_VER_NAMESPACE(_V1) diff --git a/sycl/include/sycl/properties/all_properties.hpp b/sycl/include/sycl/properties/all_properties.hpp index 91bf7f0cf584e..50a7db85a0bd1 100644 --- a/sycl/include/sycl/properties/all_properties.hpp +++ b/sycl/include/sycl/properties/all_properties.hpp @@ -1,6 +1,7 @@ #include #include #include +#include #include #include #include diff --git a/sycl/include/sycl/properties/host_task_properties.hpp b/sycl/include/sycl/properties/host_task_properties.hpp new file mode 100644 index 0000000000000..8e7015c5488b1 --- /dev/null +++ b/sycl/include/sycl/properties/host_task_properties.hpp @@ -0,0 +1,46 @@ +//==----------- buffer_properties.hpp --- SYCL buffer properties -----------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#pragma once + +#include +#include +#include + +namespace sycl { +__SYCL_INLINE_VER_NAMESPACE(_V1) { + +namespace property { +namespace host_task { +class exec_on_submit + : public detail::DataLessProperty {}; + +class manual_interop_sync + : public detail::DataLessProperty {}; + +} // namespace host_task +} // namespace property + +// Forward declaration +class host_task; + +template <> +struct is_property : std::true_type {}; +template <> +struct is_property : std::true_type { +}; + +template <> +struct is_property_of + : std::true_type {}; +template <> +struct is_property_of + : std::true_type {}; + +} // __SYCL_INLINE_VER_NAMESPACE(_V1) +} // namespace sycl diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp index 8f5f18a090aa8..e1afc48fb1478 100644 --- a/sycl/source/detail/scheduler/commands.cpp +++ b/sycl/source/detail/scheduler/commands.cpp @@ -27,6 +27,7 @@ #include #include #include +#include #include #include @@ -293,16 +294,19 @@ class DispatchHostTask { CGHostTask &HostTask = static_cast(MThisCmd->getCG()); - pi_result WaitResult = waitForEvents(); - if (WaitResult != PI_SUCCESS) { - std::exception_ptr EPtr = std::make_exception_ptr(sycl::runtime_error( - std::string("Couldn't wait for host-task's dependencies"), - WaitResult)); - HostTask.MQueue->reportAsyncException(EPtr); - - // reset host-task's lambda and quit - HostTask.MHostTask.reset(); - return; + if (!HostTask.MHostTask->MPropertyList + .has_property()) { + pi_result WaitResult = waitForEvents(); + if (WaitResult != PI_SUCCESS) { + std::exception_ptr EPtr = std::make_exception_ptr(sycl::runtime_error( + std::string("Couldn't wait for host-task's dependencies"), + WaitResult)); + HostTask.MQueue->reportAsyncException(EPtr); + + // reset host-task's lambda and quit + HostTask.MHostTask.reset(); + return; + } } try { @@ -310,7 +314,9 @@ class DispatchHostTask { if (HostTask.MHostTask->isInteropTask()) { interop_handle IH{MReqToMem, HostTask.MQueue, HostTask.MQueue->getDeviceImplPtr(), - HostTask.MQueue->getContextImplPtr()}; + HostTask.MQueue->getContextImplPtr(), + MThisCmd->getPiEvents(MThisCmd->MPreparedDepsEvents), + HostTask.MHostTask->MPropertyList}; HostTask.MHostTask->call(IH); } else From 311e07477113626d610b53e3f40e03f58d0978d7 Mon Sep 17 00:00:00 2001 From: Hugh Delaney Date: Wed, 5 Oct 2022 16:07:57 +0100 Subject: [PATCH 2/5] WIP --- sycl/include/sycl/interop_handle.hpp | 5 ----- 1 file changed, 5 deletions(-) diff --git a/sycl/include/sycl/interop_handle.hpp b/sycl/include/sycl/interop_handle.hpp index 5c1568a1f044c..2cd22a60015ff 100644 --- a/sycl/include/sycl/interop_handle.hpp +++ b/sycl/include/sycl/interop_handle.hpp @@ -143,7 +143,6 @@ class interop_handle { #endif } - template std::vector get_native_events() { #ifndef __SYCL_DEVICE_ONLY__ if (!MPropertyList @@ -152,10 +151,6 @@ class interop_handle { "get_native_events can only be used in host task " "with manual_interop_sync property"); } - if (Backend != get_backend()) { - throw sycl::exception(make_error_code(errc::backend_mismatch), - "Incorrect backend argument was passed"); - } return MRawEvents; #endif throw sycl::exception(make_error_code(errc::feature_not_supported), From 246a30b8653785eff32855b380b7b901061b32fe Mon Sep 17 00:00:00 2001 From: Hugh Delaney Date: Tue, 11 Oct 2022 10:26:34 +0100 Subject: [PATCH 3/5] WIP ih.get_native_events working --- sycl/include/sycl/detail/cg.hpp | 8 +++- sycl/include/sycl/detail/cg_types.hpp | 26 ++++++++--- sycl/include/sycl/event.hpp | 1 + sycl/include/sycl/handler.hpp | 4 +- sycl/include/sycl/interop_handle.hpp | 29 ++++++++---- sycl/source/detail/scheduler/commands.cpp | 56 +++++++++++++---------- 6 files changed, 82 insertions(+), 42 deletions(-) diff --git a/sycl/include/sycl/detail/cg.hpp b/sycl/include/sycl/detail/cg.hpp index 856b2b187e1a6..a6f482ccd4708 100644 --- a/sycl/include/sycl/detail/cg.hpp +++ b/sycl/include/sycl/detail/cg.hpp @@ -374,7 +374,13 @@ class CGHostTask : public CG { std::move(SharedPtrStorage), std::move(Requirements), std::move(Events), std::move(loc)), MHostTask(std::move(HostTask)), MQueue(Queue), MContext(Context), - MArgs(std::move(Args)) {} + MArgs(std::move(Args)) { + std::cout + << "CGHostTask(...) with HostTask has property: " + << MHostTask->MPropertyList + ->has_property() + << std::endl; + } }; class CGBarrier : public CG { diff --git a/sycl/include/sycl/detail/cg_types.hpp b/sycl/include/sycl/detail/cg_types.hpp index 5bc6a9326feb6..115e00685fd1a 100644 --- a/sycl/include/sycl/detail/cg_types.hpp +++ b/sycl/include/sycl/detail/cg_types.hpp @@ -228,20 +228,32 @@ class InteropTask { class HostTask { std::function MHostTask; std::function MInteropTask; - property_list MPropertyList; + std::shared_ptr MPropertyList; public: - HostTask() : MHostTask([]() {}) {} - HostTask(std::function &&Func) : MHostTask(Func) {} - HostTask(std::function &&Func) : MInteropTask(Func) {} + HostTask() : MHostTask([]() {}) { std::cout << "HostTask()\n"; } + HostTask(std::function &&Func, std::shared_ptr(PL)) + : MHostTask(Func), + MPropertyList(PL) { + std::cout << "HostTask(Func(), PL)\n"; + } + HostTask(std::function &&Func, + std::shared_ptr(PL)) + : MInteropTask(Func), MPropertyList(PL) { + std::cout << "HostTask(func(ih), PL)\n"; + } - bool isInteropTask() const { return !!MInteropTask; } + bool isInteropTask() const { + return !!MInteropTask; } - void call() { MHostTask(); } - void call(interop_handle handle) { MInteropTask(handle); } + void call() { + MHostTask(); } + void call(interop_handle handle) { + MInteropTask(handle); } friend class sycl::handler; friend class DispatchHostTask; + friend class CGHostTask; }; // Class which stores specific lambda object. diff --git a/sycl/include/sycl/event.hpp b/sycl/include/sycl/event.hpp index b672e1c114280..d0bb07459b486 100644 --- a/sycl/include/sycl/event.hpp +++ b/sycl/include/sycl/event.hpp @@ -147,6 +147,7 @@ class __SYCL_EXPORT event { template friend auto get_native(const SyclObjectT &Obj) -> backend_return_t; + friend class interop_handle; }; } // __SYCL_INLINE_VER_NAMESPACE(_V1) diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 9eec053224aea..21602f7e74d3b 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -1295,8 +1295,8 @@ class __SYCL_EXPORT handler { MNDRDesc.set(range<1>(1)); MArgs = std::move(MAssociatedAccesors); - MHostTask.reset(new detail::HostTask(std::move(Func))); - MHostTask->MPropertyList = std::move(PropList); + MHostTask.reset(new detail::HostTask( + std::move(Func), std::make_shared(PropList))); setType(detail::CG::CodeplayHostTask); } diff --git a/sycl/include/sycl/interop_handle.hpp b/sycl/include/sycl/interop_handle.hpp index 2cd22a60015ff..abf475cf17904 100644 --- a/sycl/include/sycl/interop_handle.hpp +++ b/sycl/include/sycl/interop_handle.hpp @@ -28,6 +28,8 @@ class DispatchHostTask; class queue_impl; class device_impl; class context_impl; + +using EventImplPtr = std::shared_ptr; } // namespace detail class queue; @@ -143,15 +145,24 @@ class interop_handle { #endif } - std::vector get_native_events() { + template + std::vector> get_native_events() { #ifndef __SYCL_DEVICE_ONLY__ if (!MPropertyList - .has_property()) { + ->has_property()) { throw sycl::exception(make_error_code(errc::feature_not_supported), "get_native_events can only be used in host task " "with manual_interop_sync property"); } - return MRawEvents; + std::vector> native_events; + for (auto &EventImplPtr : MEventImplPtrs){ + if (EventImplPtr) { + native_events.push_back(static_cast>( + get_native(event(EventImplPtr)))); + } + } + std::cout << "native_events.size(): " << native_events.size() << std::endl; + return native_events; #endif throw sycl::exception(make_error_code(errc::feature_not_supported), "get_native_events should never be called on device"); @@ -166,11 +177,11 @@ class interop_handle { const std::shared_ptr &Queue, const std::shared_ptr &Device, const std::shared_ptr &Context, - std::vector RawEvents, - property_list PropList = {}) + std::vector EventImplPtrs, + std::shared_ptr PropList = {}) : MQueue(Queue), MDevice(Device), MContext(Context), - MMemObjs(std::move(MemObjs)), MRawEvents(RawEvents), - MPropertyList(PropList) {} + MPropertyList(PropList), MMemObjs(std::move(MemObjs)), + MEventImplPtrs(EventImplPtrs) {} template backend_return_t> @@ -189,10 +200,10 @@ class interop_handle { std::shared_ptr MQueue; std::shared_ptr MDevice; std::shared_ptr MContext; + std::shared_ptr MPropertyList; std::vector MMemObjs; - std::vector MRawEvents; - property_list MPropertyList; + std::vector MEventImplPtrs; }; } // __SYCL_INLINE_VER_NAMESPACE(_V1) diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp index e1afc48fb1478..a47767aed1dc3 100644 --- a/sycl/source/detail/scheduler/commands.cpp +++ b/sycl/source/detail/scheduler/commands.cpp @@ -26,8 +26,8 @@ #include #include #include -#include #include +#include #include #include @@ -253,25 +253,26 @@ class DispatchHostTask { RequiredEventsPerPlugin[&Plugin].push_back(Event); } - // wait for dependency device events - // FIXME Current implementation of waiting for events will make the thread - // 'sleep' until all of dependency events are complete. We need a bit more - // sophisticated waiting mechanism to allow to utilize this thread for any - // other available job and resume once all required events are ready. - for (auto &PluginWithEvents : RequiredEventsPerPlugin) { - std::vector RawEvents = - MThisCmd->getPiEvents(PluginWithEvents.second); - try { - PluginWithEvents.first->call(RawEvents.size(), - RawEvents.data()); - } catch (const sycl::exception &E) { - CGHostTask &HostTask = static_cast(MThisCmd->getCG()); - HostTask.MQueue->reportAsyncException(std::current_exception()); - return (pi_result)E.get_cl_code(); - } catch (...) { - CGHostTask &HostTask = static_cast(MThisCmd->getCG()); - HostTask.MQueue->reportAsyncException(std::current_exception()); - return PI_ERROR_UNKNOWN; + std::cout << "Before dynamic cast:\n"; + CGHostTask &HostTask = dynamic_cast(MThisCmd->getCG()); + std::cout << "After dynamic cast:\n"; + if (!HostTask.MHostTask->MPropertyList + ->has_property()) { + for (auto &PluginWithEvents : RequiredEventsPerPlugin) { + std::vector RawEvents = + MThisCmd->getPiEvents(PluginWithEvents.second); + try { + PluginWithEvents.first->call( + RawEvents.size(), RawEvents.data()); + std::cout << "I'm waiting here " << __FILE__ << " : " << __LINE__ + << std::endl; + } catch (const sycl::exception &E) { + HostTask.MQueue->reportAsyncException(std::current_exception()); + return (pi_result)E.get_cl_code(); + } catch (...) { + HostTask.MQueue->reportAsyncException(std::current_exception()); + return PI_ERROR_UNKNOWN; + } } } @@ -295,7 +296,7 @@ class DispatchHostTask { CGHostTask &HostTask = static_cast(MThisCmd->getCG()); if (!HostTask.MHostTask->MPropertyList - .has_property()) { + ->has_property()) { pi_result WaitResult = waitForEvents(); if (WaitResult != PI_SUCCESS) { std::exception_ptr EPtr = std::make_exception_ptr(sycl::runtime_error( @@ -312,10 +313,11 @@ class DispatchHostTask { try { // we're ready to call the user-defined lambda now if (HostTask.MHostTask->isInteropTask()) { - interop_handle IH{MReqToMem, HostTask.MQueue, + interop_handle IH{MReqToMem, + HostTask.MQueue, HostTask.MQueue->getDeviceImplPtr(), HostTask.MQueue->getContextImplPtr(), - MThisCmd->getPiEvents(MThisCmd->MPreparedDepsEvents), + MThisCmd->MPreparedDepsEvents, HostTask.MHostTask->MPropertyList}; HostTask.MHostTask->call(IH); @@ -2466,12 +2468,18 @@ pi_int32 ExecCGCommand::enqueueImp() { return PI_SUCCESS; } case CG::CGTYPE::CodeplayInteropTask: { + std::cout << "In CodeplayInteropTask " << __FILE__ << " : " << __LINE__ + << std::endl; const detail::plugin &Plugin = MQueue->getPlugin(); CGInteropTask *ExecInterop = (CGInteropTask *)MCommandGroup.get(); // Wait for dependencies to complete before dispatching work on the host // TODO: Use a callback to dispatch the interop task instead of waiting for // the event + // + // We don't want this wait to happen if (!RawEvents.empty()) { + std::cout << "We are waiting in " << __FILE__ << " : " << __LINE__ + << std::endl; Plugin.call(RawEvents.size(), &RawEvents[0]); } std::vector ReqMemObjs; @@ -2495,6 +2503,8 @@ pi_int32 ExecCGCommand::enqueueImp() { return PI_SUCCESS; } case CG::CGTYPE::CodeplayHostTask: { + std::cout << "In CodeplayHostTask " << __FILE__ << " : " << __LINE__ + << std::endl; CGHostTask *HostTask = static_cast(MCommandGroup.get()); for (ArgDesc &Arg : HostTask->MArgs) { From 6434ecaef6c4f19c10f431383d7aa9385756a0e9 Mon Sep 17 00:00:00 2001 From: Hugh Delaney Date: Tue, 18 Oct 2022 16:15:46 +0100 Subject: [PATCH 4/5] WIP --- sycl/include/sycl/detail/cg.hpp | 13 +- sycl/include/sycl/detail/cg_types.hpp | 52 +++++++- sycl/include/sycl/handler.hpp | 27 +++- sycl/include/sycl/interop_handle.hpp | 43 ++++-- sycl/out.txt | 125 ++++++++++++++++++ sycl/source/detail/event_impl.cpp | 14 ++ sycl/source/detail/event_impl.hpp | 60 +++++++++ sycl/source/detail/helpers.cpp | 1 + sycl/source/detail/queue_impl.cpp | 2 + sycl/source/detail/scheduler/commands.cpp | 102 +++++++++----- sycl/source/detail/scheduler/commands.hpp | 3 +- .../source/detail/scheduler/graph_builder.cpp | 8 +- sycl/source/detail/thread_pool.hpp | 17 +++ sycl/source/handler.cpp | 11 ++ 14 files changed, 420 insertions(+), 58 deletions(-) create mode 100644 sycl/out.txt diff --git a/sycl/include/sycl/detail/cg.hpp b/sycl/include/sycl/detail/cg.hpp index a6f482ccd4708..3f1c3f8b1d38f 100644 --- a/sycl/include/sycl/detail/cg.hpp +++ b/sycl/include/sycl/detail/cg.hpp @@ -375,11 +375,14 @@ class CGHostTask : public CG { std::move(Events), std::move(loc)), MHostTask(std::move(HostTask)), MQueue(Queue), MContext(Context), MArgs(std::move(Args)) { - std::cout - << "CGHostTask(...) with HostTask has property: " - << MHostTask->MPropertyList - ->has_property() - << std::endl; + if (MHostTask) { + std::cout << "CGHostTask(...) with HostTask has property: " << std::flush; + std::cout << MHostTask->hasProperty< + sycl::property::host_task::manual_interop_sync>() + << std::endl; + std::cout << "MHostTask->has_native_events(): " + << MHostTask->hasNativeEvents() << std::endl; + } } }; diff --git a/sycl/include/sycl/detail/cg_types.hpp b/sycl/include/sycl/detail/cg_types.hpp index 115e00685fd1a..082b1b2025ef0 100644 --- a/sycl/include/sycl/detail/cg_types.hpp +++ b/sycl/include/sycl/detail/cg_types.hpp @@ -242,20 +242,62 @@ class HostTask { : MInteropTask(Func), MPropertyList(PL) { std::cout << "HostTask(func(ih), PL)\n"; } + HostTask(std::shared_ptr(PL)) + : MPropertyList(PL) { + std::cout << "HostTask(PL)\n"; + } + + template + bool hasProperty() { + if (MPropertyList) { + return MPropertyList->has_property(); + } + return false; + } - bool isInteropTask() const { + virtual bool isInteropTask() const { return !!MInteropTask; } - void call() { - MHostTask(); } - void call(interop_handle handle) { - MInteropTask(handle); } + void call() { MHostTask(); } + virtual void call(interop_handle handle) { MInteropTask(handle); } + virtual bool hasNativeEvents() const { return false; } friend class sycl::handler; friend class DispatchHostTask; friend class CGHostTask; }; +template class NativeEventsHostTask : public HostTask { + std::function(interop_handle)> MHostTask; + backend_return_t MNativeEvents; + +public: + NativeEventsHostTask() = delete; + NativeEventsHostTask( + std::function(interop_handle)> &&Func, + std::shared_ptr PL) + : MHostTask(Func), HostTask(PL) { + std::cout << "NativeEventsHostTask(std::function(interop_handle)> &&Func)" + << std::endl; + std::cout << "PL.has_property: " + << hasProperty() + << std::endl; + }; + + void call(interop_handle handle) override { + MNativeEvents = MHostTask(handle); + std::cout << "Got native events from kernel: " << MNativeEvents.size() + << std::endl; + } + bool hasNativeEvents() const override { return true; } + backend_return_t getNativeEvents() const { + return MNativeEvents; + } + bool isInteropTask() const override { return true; } + backend getBackend() const { return Backend; }; +}; + // Class which stores specific lambda object. template class HostKernel : public HostKernelBase { diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 21602f7e74d3b..6a9386fb101c2 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -1300,6 +1300,24 @@ class __SYCL_EXPORT handler { setType(detail::CG::CodeplayHostTask); } + + template + detail::enable_if_t, + backend_return_t(interop_handle)>::value> + host_task_impl(FuncT &&Func, const property_list &PropList) { + throwIfActionIsCreated(); + + std::cout << "Calling this special one!\n"; + + MNDRDesc.set(range<1>(1)); + MArgs = std::move(MAssociatedAccesors); + + MHostTask.reset(new detail::NativeEventsHostTask( + std::move(Func), std::make_shared(PropList))); + + setType(detail::CG::CodeplayHostTask); + } public: handler(const handler &) = delete; @@ -1496,7 +1514,7 @@ class __SYCL_EXPORT handler { } /// Enqueues a command to the SYCL runtime to invoke \p Func once. - template + template detail::enable_if_t< detail::check_fn_signature, void()>::value || @@ -1506,6 +1524,13 @@ class __SYCL_EXPORT handler { host_task_impl(Func, PropList); } + template + detail::enable_if_t, + backend_return_t(interop_handle)>::value> + host_task(FuncT &&Func, const property_list &PropList = {}) { + host_task_impl(Func, PropList); + } // replace _KERNELFUNCPARAM(KernelFunc) with KernelType KernelFunc // or const KernelType &KernelFunc #ifdef __SYCL_NONCONST_FUNCTOR__ diff --git a/sycl/include/sycl/interop_handle.hpp b/sycl/include/sycl/interop_handle.hpp index abf475cf17904..d26bacf7465a6 100644 --- a/sycl/include/sycl/interop_handle.hpp +++ b/sycl/include/sycl/interop_handle.hpp @@ -30,6 +30,11 @@ class device_impl; class context_impl; using EventImplPtr = std::shared_ptr; + +extern bool hasNativeInteropEvents(EventImplPtr); +template +extern backend_return_t getNativeInteropEvents(EventImplPtr); +extern void waitOnNativeInteropEvents(EventImplPtr); } // namespace detail class queue; @@ -146,19 +151,39 @@ class interop_handle { } template - std::vector> get_native_events() { + backend_return_t get_native_events() { #ifndef __SYCL_DEVICE_ONLY__ if (!MPropertyList - ->has_property()) { + ->has_property()) { throw sycl::exception(make_error_code(errc::feature_not_supported), "get_native_events can only be used in host task " "with manual_interop_sync property"); } - std::vector> native_events; - for (auto &EventImplPtr : MEventImplPtrs){ - if (EventImplPtr) { - native_events.push_back(static_cast>( - get_native(event(EventImplPtr)))); + if (Backend != backend::opencl) { + throw sycl::exception(make_error_code(errc::feature_not_supported), + "Get native events is only supported in openCL "); + } + std::cout << "ih.get_native_events(): MEventImplPtrs.size() = " + << MDepEventImplPtrs.size() << std::endl; + backend_return_t native_events; + for (auto &eventImplPtr : MDepEventImplPtrs) { + if (eventImplPtr) { + auto eventImplPtrNativeEvents = + get_native(event(eventImplPtr)); + native_events.insert(native_events.end(), + eventImplPtrNativeEvents.begin(), + eventImplPtrNativeEvents.end()); + + // If previously enqueued host_tasks have returned native events, + // add them to the vector of events + if (hasNativeInteropEvents(eventImplPtr)) { + std::cout << "In here!!12345\n"; + waitOnNativeInteropEvents(eventImplPtr); + auto nativeInteropEvents = + detail::getNativeInteropEvents(eventImplPtr); + native_events.insert(native_events.end(), nativeInteropEvents.begin(), + nativeInteropEvents.end()); + } } } std::cout << "native_events.size(): " << native_events.size() << std::endl; @@ -181,7 +206,7 @@ class interop_handle { std::shared_ptr PropList = {}) : MQueue(Queue), MDevice(Device), MContext(Context), MPropertyList(PropList), MMemObjs(std::move(MemObjs)), - MEventImplPtrs(EventImplPtrs) {} + MDepEventImplPtrs(EventImplPtrs) {} template backend_return_t> @@ -203,7 +228,7 @@ class interop_handle { std::shared_ptr MPropertyList; std::vector MMemObjs; - std::vector MEventImplPtrs; + std::vector MDepEventImplPtrs; }; } // __SYCL_INLINE_VER_NAMESPACE(_V1) diff --git a/sycl/out.txt b/sycl/out.txt new file mode 100644 index 0000000000000..38479bb5a49fa --- /dev/null +++ b/sycl/out.txt @@ -0,0 +1,125 @@ +MIsNotFinalized +Command::addDep +Command::addDep +Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 +Command::addDep +Command::addDep +Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 +Command::addDep +Command::addDep +Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 +Command::addDep +Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 +Command::addDep +Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 +Command::addDep +Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 +Command::addDep +Command::addDep +Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 +HostTask(Func(), PL) +MIsNotFinalized +CGHostTask(...) with HostTask has property: 0 +MHostTask->has_native_events(): 0 +Command::addDep +Command::addDep +Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 +Command::addDep +Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 +Command::addDep +The contexts match so we are pushing back a DepEvent! +Command::addDep +Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 +Command::addDep +The contexts match so we are pushing back a DepEvent! +Command::addDep +Command::addDep +Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 +Command::addDep +Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 +Command::addDep +The contexts match so we are pushing back a DepEvent! +Command::addDep +Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 +Command::addDep +The contexts match so we are pushing back a DepEvent! +Command::addDep +Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 +Command::addDep +Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 +Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp226 +Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp226 +Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp226 +Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp226 +In CodeplayHostTask /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp2527 : 2527 +setHasNativeEvents(0); +Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp255 +Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp255 +Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp226 +Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp226 +Events.push_back: /home/hugh/llvm/sycl/source/detail/queue_impl.cpp212 +Events.push_back: /home/hugh/llvm/sycl/source/detail/queue_impl.cpp212 +Command has deps: 2 + + +Calling this special one! +HostTask(PL) +NativeEventsHostTask(std::function(interop_handle)> &&Func) +PL.has_property: 1 +MIsNotFinalized +CGHostTask(...) with HostTask has property: 1 +MHostTask->has_native_events(): 1 +Command::addDep +The contexts match so we are pushing back a DepEvent! +Command::addDep +Command::addDep +Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 +Command::addDep +Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 +Command::addDep +Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 +Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp226 +In CodeplayHostTask /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp2527 : 2527 +setHasNativeEvents(1); +Events.push_back: /home/hugh/llvm/sycl/source/detail/queue_impl.cpp212 +sycl_e2.get_wait_list: Making interop_handle with dep events: 1 +1 +ih.get_native_events(): MEventImplPtrs.size() = 1 +hasNativeInteropEvents: 0 +native_events.size(): 1 +Have 1 native events in host_task 1! +Native event at address: 0x2325348 +Returning native event at address: 0x7fbc00002498 +Got native events from kernel: 1 +Getting native event dependencies for new host_task! +In setNativeEvents with native_events.size: 1 +Command has deps: 1 + + +Calling this special one! +HostTask(PL) +NativeEventsHostTask(std::function(interop_handle)> &&Func) +PL.has_property: 1 +MIsNotFinalized +CGHostTask(...) with HostTask has property: 1 +MHostTask->has_native_events(): 1 +Command::addDep +Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 +Command::addDep +Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 +Command::addDep +Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 +In CodeplayHostTask /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp2527 : 2527 +setHasNativeEvents(1); +Events.push_back: /home/hugh/llvm/sycl/source/detail/queue_impl.cpp212 +Command::addDep +Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 +Making interop_handle with dep events: 0 +ih.get_native_events(): MEventImplPtrs.size() = 0 +native_events.size(): 0 +Have 0 native events in host_task 2! +Got native events from kernel: 1 +Getting native event dependencies for new host_task! +In setNativeEvents with native_events.size: 1 +Command has deps: 2 +success diff --git a/sycl/source/detail/event_impl.cpp b/sycl/source/detail/event_impl.cpp index 7d88e52baf27d..b843c31c2b36e 100644 --- a/sycl/source/detail/event_impl.cpp +++ b/sycl/source/detail/event_impl.cpp @@ -445,4 +445,18 @@ void event_impl::cleanDepEventsThroughOneLevel() { } // namespace detail } // __SYCL_INLINE_VER_NAMESPACE(_V1) +__SYCL_EXPORT bool +detail::hasNativeInteropEvents(std::shared_ptr EPtr) { + std::cout << "hasNativeInteropEvents: " << EPtr->hasNativeEvents() << std::endl; + return EPtr->hasNativeEvents(); +} +__SYCL_EXPORT void +detail::waitOnNativeInteropEvents(std::shared_ptr EPtr) { + return EPtr->waitOnNativeInteropEvents(); +} + +template __SYCL_EXPORT backend_return_t + detail::getNativeInteropEvents( + std::shared_ptr); + } // namespace sycl diff --git a/sycl/source/detail/event_impl.hpp b/sycl/source/detail/event_impl.hpp index d33da055a4f8e..00d62426d9dd6 100644 --- a/sycl/source/detail/event_impl.hpp +++ b/sycl/source/detail/event_impl.hpp @@ -9,6 +9,7 @@ #pragma once #include +#include #include #include #include @@ -33,6 +34,22 @@ using QueueImplPtr = std::shared_ptr; class event_impl; using EventImplPtr = std::shared_ptr; +class nativeEventsBuffer { +public: + virtual ~nativeEventsBuffer() = default; +}; + +template +class backendNativeEventsBuffer : public nativeEventsBuffer { + backend_return_t nativeEvents; + +public: + backendNativeEventsBuffer() = delete; + backendNativeEventsBuffer(backend_return_t Evs) + : nativeEvents(Evs){}; + backend_return_t& getEvents() { return nativeEvents; } +}; + class event_impl { public: enum HostEventState : int { @@ -172,6 +189,38 @@ class event_impl { /// \return a native handle. pi_native_handle getNative(); + template + void setNativeInteropEvents(backend_return_t nativeEvents) { + assert(!MNativeEvents && "setNativeEvents should only be called once per" + "event_impl"); + std::cout << "In setNativeEvents with native_events.size: " + << nativeEvents.size() << "\n"; + MNativeEvents.reset(new backendNativeEventsBuffer(nativeEvents)); + { + auto ul = std::unique_lock(MMutex); + MNativeEventsPopulated = true; + MIsInitialized = true; + } + cv.notify_one(); + } + + template + backend_return_t getNativeInteropEvents() { + auto evs = + dynamic_cast *>(MNativeEvents.get()) + ->getEvents(); + std::cout << "getNativeInteropEvents has size: " << evs.size() << std::endl; + return evs; + } + + bool hasNativeEvents() const { return MHasNativeEvents; } + void setHasNativeEvents(const bool val) { MHasNativeEvents = val; } + void waitOnNativeInteropEvents() { + auto ul = std::unique_lock(MMutex); + cv.wait(ul, [this]() { return MNativeEventsPopulated; }); + std::cout << "nativeEvents has been populated!\n"; + } + /// Returns vector of event dependencies. /// /// @return a reference to MPreparedDepsEvents. @@ -266,6 +315,10 @@ class event_impl { /// the queue to the device. std::atomic MIsFlushed = false; + bool MHasNativeEvents = false; + bool MNativeEventsPopulated = false; + std::unique_ptr MNativeEvents; + // State of host event. Employed only for host events and event with no // backend's representation (e.g. alloca). Used values are listed in // HostEventState enum. @@ -280,6 +333,8 @@ class event_impl { std::mutex MMutex; std::condition_variable cv; + friend class interop_handle; + friend std::vector getOrWaitEvents(std::vector DepEvents, std::shared_ptr Context); @@ -287,4 +342,9 @@ class event_impl { } // namespace detail } // __SYCL_INLINE_VER_NAMESPACE(_V1) +template +backend_return_t +detail::getNativeInteropEvents(std::shared_ptr EPtr) { + return EPtr->getNativeInteropEvents(); +} } // namespace sycl diff --git a/sycl/source/detail/helpers.cpp b/sycl/source/detail/helpers.cpp index c605b845b0147..225206e59970e 100644 --- a/sycl/source/detail/helpers.cpp +++ b/sycl/source/detail/helpers.cpp @@ -34,6 +34,7 @@ std::vector getOrWaitEvents(std::vector DepEvents, SyclEventImplPtr->getContextImpl() != Context) { SyclEventImplPtr->waitInternal(); } else { + std::cout << "Events.push_back: " << __FILE__ << std::endl; Events.push_back(SyclEventImplPtr->getHandleRef()); } } diff --git a/sycl/source/detail/queue_impl.cpp b/sycl/source/detail/queue_impl.cpp index 9b88f83789b16..79c785e654281 100644 --- a/sycl/source/detail/queue_impl.cpp +++ b/sycl/source/detail/queue_impl.cpp @@ -209,6 +209,7 @@ void queue_impl::addEvent(const event &Event) { std::weak_ptr EventWeakPtr{EImpl}; std::lock_guard Lock{MMutex}; MEventsWeak.push_back(std::move(EventWeakPtr)); + std::cout << "Events.push_back: " << __FILE__ << __LINE__ << std::endl; } } @@ -241,6 +242,7 @@ void queue_impl::addSharedEvent(const event &Event) { info::event_command_status::complete; })); } + std::cout << "Events.push_back: " << __FILE__ << __LINE__<< std::endl; MEventsShared.push_back(Event); } diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp index a47767aed1dc3..942ca22586f69 100644 --- a/sycl/source/detail/scheduler/commands.cpp +++ b/sycl/source/detail/scheduler/commands.cpp @@ -23,6 +23,7 @@ #include #include #include +#include #include #include #include @@ -222,6 +223,7 @@ Command::getPiEvents(const std::vector &EventImpls) const { continue; RetPiEvents.push_back(EventImpl->getHandleRef()); + std::cout << "Events.push_back: " << __FILE__ << __LINE__ << std::endl; } return RetPiEvents; @@ -250,29 +252,23 @@ class DispatchHostTask { for (const EventImplPtr &Event : MThisCmd->MPreparedDepsEvents) { const detail::plugin &Plugin = Event->getPlugin(); + std::cout << "Events.push_back: " << __FILE__ << __LINE__ << std::endl; RequiredEventsPerPlugin[&Plugin].push_back(Event); } - std::cout << "Before dynamic cast:\n"; CGHostTask &HostTask = dynamic_cast(MThisCmd->getCG()); - std::cout << "After dynamic cast:\n"; - if (!HostTask.MHostTask->MPropertyList - ->has_property()) { - for (auto &PluginWithEvents : RequiredEventsPerPlugin) { - std::vector RawEvents = - MThisCmd->getPiEvents(PluginWithEvents.second); - try { - PluginWithEvents.first->call( - RawEvents.size(), RawEvents.data()); - std::cout << "I'm waiting here " << __FILE__ << " : " << __LINE__ - << std::endl; - } catch (const sycl::exception &E) { - HostTask.MQueue->reportAsyncException(std::current_exception()); - return (pi_result)E.get_cl_code(); - } catch (...) { - HostTask.MQueue->reportAsyncException(std::current_exception()); - return PI_ERROR_UNKNOWN; - } + for (auto &PluginWithEvents : RequiredEventsPerPlugin) { + std::vector RawEvents = + MThisCmd->getPiEvents(PluginWithEvents.second); + try { + PluginWithEvents.first->call(RawEvents.size(), + RawEvents.data()); + } catch (const sycl::exception &E) { + HostTask.MQueue->reportAsyncException(std::current_exception()); + return (pi_result)E.get_cl_code(); + } catch (...) { + HostTask.MQueue->reportAsyncException(std::current_exception()); + return PI_ERROR_UNKNOWN; } } @@ -313,6 +309,8 @@ class DispatchHostTask { try { // we're ready to call the user-defined lambda now if (HostTask.MHostTask->isInteropTask()) { + std::cout << "Making interop_handle with dep events: " + << MThisCmd->MPreparedDepsEvents.size() << "\n"; interop_handle IH{MReqToMem, HostTask.MQueue, HostTask.MQueue->getDeviceImplPtr(), @@ -327,8 +325,22 @@ class DispatchHostTask { HostTask.MQueue->reportAsyncException(std::current_exception()); } - HostTask.MHostTask.reset(); + if (HostTask.MHostTask->hasNativeEvents()) { + auto NativeEventsHT = + dynamic_cast *>( + HostTask.MHostTask.get()); + if (NativeEventsHT && NativeEventsHT->getBackend() == backend::opencl) { + std::cout << "Getting native event dependencies for new host_task!\n"; + MThisCmd->MEvent->setNativeInteropEvents( + NativeEventsHT->getNativeEvents()); + } else { + throw sycl::exception(errc::invalid, + "Host task manual interop" + "sync is currently only enabled for opencl"); + } + } + HostTask.MHostTask.reset(); // unblock user empty command here EmptyCommand *EmptyCmd = MThisCmd->MEmptyCmd; assert(EmptyCmd && "No empty command found"); @@ -345,6 +357,7 @@ class DispatchHostTask { Scheduler::ReadLockT Lock(Sched.MGraphLock); std::vector Deps = MThisCmd->MDeps; + std::cout << "Command has deps: " << MThisCmd->MDeps.size() << std::endl; // update self-event status MThisCmd->MEvent->setComplete(); @@ -390,6 +403,7 @@ void Command::waitForEvents(QueueImplPtr Queue, ContextImplPtr Context = Event->getContextImpl(); assert(Context.get() && "Only non-host events are expected to be waited for here"); + std::cout << "Events.push_back: " << __FILE__ << __LINE__ << std::endl; RequiredEventsPerContext[Context.get()].push_back(Event); } @@ -419,6 +433,7 @@ void Command::waitForEvents(QueueImplPtr Queue, /// should not outlive the event connected to it. Command::Command(CommandType Type, QueueImplPtr Queue) : MQueue(std::move(Queue)), + // HUGHTODO PROBLEM HERE MEvent(std::make_shared(MQueue)), MPreparedDepsEvents(MEvent->getPreparedDepsEvents()), MPreparedHostDepsEvents(MEvent->getPreparedHostDepsEvents()), @@ -616,13 +631,15 @@ Command *Command::processDepEvent(EventImplPtr DepEvent, const DepDesc &Dep, // (e.g. alloca). Note that we can't check the pi event to make that // distinction since the command might still be unenqueued at this point. bool PiEventExpected = (!DepEvent->is_host() && DepEvent->isInitialized()) || + DepEvent->hasNativeEvents() || getType() == CommandType::HOST_TASK; if (auto *DepCmd = static_cast(DepEvent->getCommand())) - PiEventExpected &= DepCmd->producesPiEvent(); + PiEventExpected &= DepCmd->producesPiEvent() || DepEvent->hasNativeEvents(); if (!PiEventExpected) { // call to waitInternal() is in waitForPreparedHostEvents() as it's called // from enqueue process functions + std::cout << "Events.push_back: " << __FILE__ << __LINE__ << std::endl; MPreparedHostDepsEvents.push_back(DepEvent); return nullptr; } @@ -632,11 +649,14 @@ Command *Command::processDepEvent(EventImplPtr DepEvent, const DepDesc &Dep, ContextImplPtr DepEventContext = DepEvent->getContextImpl(); // If contexts don't match we'll connect them using host task if (DepEventContext != WorkerContext && !WorkerContext->is_host()) { + std::cout << "The contexts don't match so we are connecting these things " + "with a host task!\n"; Scheduler::GraphBuilder &GB = Scheduler::getInstance().MGraphBuilder; ConnectionCmd = GB.connectDepEvent(this, DepEvent, Dep, ToCleanUp); - } else + } else { MPreparedDepsEvents.push_back(std::move(DepEvent)); - + std::cout << "The contexts match so we are pushing back a DepEvent!\n"; + } return ConnectionCmd; } @@ -656,17 +676,19 @@ bool Command::supportsPostEnqueueCleanup() const { return true; } Command *Command::addDep(DepDesc NewDep, std::vector &ToCleanUp) { Command *ConnectionCmd = nullptr; + std::cout << "Command::addDep\n"; + if (NewDep.MDepCommand) { - ConnectionCmd = + ConnectionCmd = // HT1 processDepEvent(NewDep.MDepCommand->getEvent(), NewDep, ToCleanUp); } // ConnectionCmd insertion builds the following dependency structure: // this -> emptyCmd (for ConnectionCmd) -> ConnectionCmd -> NewDep // that means that this and NewDep are already dependent if (!ConnectionCmd) { - MDeps.push_back(NewDep); + MDeps.push_back(NewDep); // HT1 if (NewDep.MDepCommand) - NewDep.MDepCommand->addUser(this); + NewDep.MDepCommand->addUser(this); // HT1 } #ifdef XPTI_ENABLE_INSTRUMENTATION @@ -2468,8 +2490,8 @@ pi_int32 ExecCGCommand::enqueueImp() { return PI_SUCCESS; } case CG::CGTYPE::CodeplayInteropTask: { - std::cout << "In CodeplayInteropTask " << __FILE__ << " : " << __LINE__ - << std::endl; + std::cout << "In CodeplayInteropTask " << __FILE__ << __LINE__ << " : " + << __LINE__ << std::endl; const detail::plugin &Plugin = MQueue->getPlugin(); CGInteropTask *ExecInterop = (CGInteropTask *)MCommandGroup.get(); // Wait for dependencies to complete before dispatching work on the host @@ -2478,8 +2500,8 @@ pi_int32 ExecCGCommand::enqueueImp() { // // We don't want this wait to happen if (!RawEvents.empty()) { - std::cout << "We are waiting in " << __FILE__ << " : " << __LINE__ - << std::endl; + std::cout << "We are waiting in " << __FILE__ << __LINE__ << " : " + << __LINE__ << std::endl; Plugin.call(RawEvents.size(), &RawEvents[0]); } std::vector ReqMemObjs; @@ -2503,8 +2525,8 @@ pi_int32 ExecCGCommand::enqueueImp() { return PI_SUCCESS; } case CG::CGTYPE::CodeplayHostTask: { - std::cout << "In CodeplayHostTask " << __FILE__ << " : " << __LINE__ - << std::endl; + std::cout << "In CodeplayHostTask " << __FILE__ << __LINE__ << " : " + << __LINE__ << std::endl; CGHostTask *HostTask = static_cast(MCommandGroup.get()); for (ArgDesc &Arg : HostTask->MArgs) { @@ -2552,8 +2574,20 @@ pi_int32 ExecCGCommand::enqueueImp() { std::sort(std::begin(ReqToMem), std::end(ReqToMem)); } - MQueue->getThreadPool().submit( - DispatchHostTask(this, std::move(ReqToMem))); + std::cout << "setHasNativeEvents(" << HostTask->MHostTask->hasNativeEvents() + << ");\n"; + this->MEvent->setHasNativeEvents(HostTask->MHostTask->hasNativeEvents()); + + if (HostTask->MHostTask + ->hasProperty()) { + std::cout << "Waiting on submit operation to complete\n"; + MQueue->getThreadPool().submit_and_wait( + DispatchHostTask(this, std::move(ReqToMem))); + + } else { + MQueue->getThreadPool().submit( + DispatchHostTask(this, std::move(ReqToMem))); + } MShouldCompleteEventIfPossible = false; diff --git a/sycl/source/detail/scheduler/commands.hpp b/sycl/source/detail/scheduler/commands.hpp index 77afa4936bc0a..d35c0e12289f2 100644 --- a/sycl/source/detail/scheduler/commands.hpp +++ b/sycl/source/detail/scheduler/commands.hpp @@ -189,7 +189,8 @@ class Command { return nullptr; } - virtual ~Command() { MEvent->cleanDepEventsThroughOneLevel(); } + virtual ~Command() { + MEvent->cleanDepEventsThroughOneLevel(); } const char *getBlockReason() const; diff --git a/sycl/source/detail/scheduler/graph_builder.cpp b/sycl/source/detail/scheduler/graph_builder.cpp index e4a609566d4cc..f6dbf339bb2bf 100644 --- a/sycl/source/detail/scheduler/graph_builder.cpp +++ b/sycl/source/detail/scheduler/graph_builder.cpp @@ -565,12 +565,15 @@ Scheduler::GraphBuilder::findDepsForReq(MemObjRecord *Record, std::vector Visited; const bool ReadOnlyReq = Req->MAccessMode == access::mode::read; + // Size 1 for host_task 1 std::vector ToAnalyze{Record->MWriteLeaves.toVector()}; if (!ReadOnlyReq) { + // HT1 in here std::vector V{Record->MReadLeaves.toVector()}; - ToAnalyze.insert(ToAnalyze.begin(), V.begin(), V.end()); + ToAnalyze.insert(ToAnalyze.begin(), V.begin(), + V.end()); // HT1 nothing inserted } while (!ToAnalyze.empty()) { @@ -579,7 +582,7 @@ Scheduler::GraphBuilder::findDepsForReq(MemObjRecord *Record, std::vector NewAnalyze; - for (const DepDesc &Dep : DepCmd->MDeps) { + for (const DepDesc &Dep : DepCmd->MDeps) { // HT1 size 6 if (Dep.MDepRequirement->MSYCLMemObj != Req->MSYCLMemObj) continue; @@ -1023,7 +1026,6 @@ Scheduler::GraphBuilder::addCG(std::unique_ptr CommandGroup, NewCmd->MEmptyCmd = addEmptyCmd(NewCmd.get(), NewCmd->getCG().MRequirements, Queue, Command::BlockReason::HostTask, ToEnqueue); - if (MPrintOptionsArray[AfterAddCG]) printGraphAsDot("after_addCG"); diff --git a/sycl/source/detail/thread_pool.hpp b/sycl/source/detail/thread_pool.hpp index 2ffbe8a0bd52a..1051a5b2db2b6 100644 --- a/sycl/source/detail/thread_pool.hpp +++ b/sycl/source/detail/thread_pool.hpp @@ -85,6 +85,23 @@ class ThreadPool { MDoSmthOrStop.notify_one(); } + + template void submit_and_wait(T &&Func) { + std::atomic_bool job_finished(false); + { + std::lock_guard Lock(MJobQueueMutex); + MJobQueue.emplace([F = std::move(Func), &job_finished]() { + F(); + job_finished.store(true); + }); + } + + MDoSmthOrStop.notify_one(); + + // Wait until job has finished + while (!job_finished.load()) + ; + } void submit(std::function &&Func) { { diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index bc3aba5d80fbc..54124047882be 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -86,12 +86,14 @@ void handler::setHandlerKernelBundle(kernel Kernel) { setHandlerKernelBundle(KernelBundleImpl); } +// HUGHTODO this is where I need to work event handler::finalize() { // This block of code is needed only for reduction implementation. // It is harmless (does nothing) for everything else. if (MIsFinalized) return MLastEvent; MIsFinalized = true; + std::cout << "MIsNotFinalized\n"; const auto &type = getType(); if (type == detail::CG::Kernel) { @@ -100,8 +102,10 @@ event handler::finalize() { getOrInsertHandlerKernelBundle(/*Insert=*/false); if (KernelBundleImpPtr) { // Make sure implicit non-interop kernel bundles have the kernel + std::cout << "KernelBundleImplPtr\n"; if (!KernelBundleImpPtr->isInterop() && !MImpl->isStateExplicitKernelBundle()) { + std::cout << "!KernelBundleImplPtr->IsInterop\n"; kernel_id KernelID = detail::ProgramManager::getInstance().getSYCLKernelID(MKernelName); bool KernelInserted = @@ -151,6 +155,7 @@ event handler::finalize() { // the graph is not changed, then this faster path is used to submit // kernel bypassing scheduler and avoiding CommandGroup, Command objects // creation. + std::cout << __LINE__ << "\n"; std::vector RawEvents; detail::EventImplPtr NewEvent; @@ -185,6 +190,7 @@ event handler::finalize() { bool DiscardEvent = false; if (MQueue->has_discard_events_support()) { + std::cout << __LINE__ << "\n"; // Kernel only uses assert if it's non interop one bool KernelUsesAssert = !(MKernel && MKernel->isInterop()) && @@ -194,10 +200,12 @@ event handler::finalize() { } if (DiscardEvent) { + std::cout << __LINE__ << "\n"; if (PI_SUCCESS != EnqueueKernel()) throw runtime_error("Enqueue process failed.", PI_ERROR_INVALID_OPERATION); } else { + std::cout << __LINE__ << "\n"; NewEvent = std::make_shared(MQueue); NewEvent->setContextImpl(MQueue->getContextImplPtr()); NewEvent->setStateIncomplete(); @@ -232,6 +240,7 @@ event handler::finalize() { break; } case detail::CG::CodeplayInteropTask: + std::cout << __LINE__ << "\n"; CommandGroup.reset(new detail::CGInteropTask( std::move(MInteropTask), std::move(MArgsStorage), std::move(MAccStorage), std::move(MSharedPtrStorage), @@ -688,6 +697,7 @@ void handler::depends_on(event Event) { "Queue operation cannot depend on discarded event."); } MEvents.push_back(EventImpl); + std::cout << "Events.push_back: " << __FILE__ << std::endl; } void handler::depends_on(const std::vector &Events) { @@ -698,6 +708,7 @@ void handler::depends_on(const std::vector &Events) { make_error_code(errc::invalid), "Queue operation cannot depend on discarded event."); } + std::cout << "Events.push_back: " << __FILE__ << std::endl; MEvents.push_back(EventImpl); } } From 6cd9db20dc02fef9f6fb4a0678df255744a65805 Mon Sep 17 00:00:00 2001 From: Hugh Delaney <46290137+hdelan@users.noreply.github.com> Date: Thu, 23 Nov 2023 17:00:14 +0000 Subject: [PATCH 5/5] Delete sycl/out.txt --- sycl/out.txt | 125 --------------------------------------------------- 1 file changed, 125 deletions(-) delete mode 100644 sycl/out.txt diff --git a/sycl/out.txt b/sycl/out.txt deleted file mode 100644 index 38479bb5a49fa..0000000000000 --- a/sycl/out.txt +++ /dev/null @@ -1,125 +0,0 @@ -MIsNotFinalized -Command::addDep -Command::addDep -Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 -Command::addDep -Command::addDep -Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 -Command::addDep -Command::addDep -Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 -Command::addDep -Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 -Command::addDep -Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 -Command::addDep -Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 -Command::addDep -Command::addDep -Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 -HostTask(Func(), PL) -MIsNotFinalized -CGHostTask(...) with HostTask has property: 0 -MHostTask->has_native_events(): 0 -Command::addDep -Command::addDep -Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 -Command::addDep -Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 -Command::addDep -The contexts match so we are pushing back a DepEvent! -Command::addDep -Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 -Command::addDep -The contexts match so we are pushing back a DepEvent! -Command::addDep -Command::addDep -Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 -Command::addDep -Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 -Command::addDep -The contexts match so we are pushing back a DepEvent! -Command::addDep -Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 -Command::addDep -The contexts match so we are pushing back a DepEvent! -Command::addDep -Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 -Command::addDep -Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 -Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp226 -Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp226 -Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp226 -Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp226 -In CodeplayHostTask /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp2527 : 2527 -setHasNativeEvents(0); -Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp255 -Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp255 -Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp226 -Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp226 -Events.push_back: /home/hugh/llvm/sycl/source/detail/queue_impl.cpp212 -Events.push_back: /home/hugh/llvm/sycl/source/detail/queue_impl.cpp212 -Command has deps: 2 - - -Calling this special one! -HostTask(PL) -NativeEventsHostTask(std::function(interop_handle)> &&Func) -PL.has_property: 1 -MIsNotFinalized -CGHostTask(...) with HostTask has property: 1 -MHostTask->has_native_events(): 1 -Command::addDep -The contexts match so we are pushing back a DepEvent! -Command::addDep -Command::addDep -Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 -Command::addDep -Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 -Command::addDep -Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 -Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp226 -In CodeplayHostTask /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp2527 : 2527 -setHasNativeEvents(1); -Events.push_back: /home/hugh/llvm/sycl/source/detail/queue_impl.cpp212 -sycl_e2.get_wait_list: Making interop_handle with dep events: 1 -1 -ih.get_native_events(): MEventImplPtrs.size() = 1 -hasNativeInteropEvents: 0 -native_events.size(): 1 -Have 1 native events in host_task 1! -Native event at address: 0x2325348 -Returning native event at address: 0x7fbc00002498 -Got native events from kernel: 1 -Getting native event dependencies for new host_task! -In setNativeEvents with native_events.size: 1 -Command has deps: 1 - - -Calling this special one! -HostTask(PL) -NativeEventsHostTask(std::function(interop_handle)> &&Func) -PL.has_property: 1 -MIsNotFinalized -CGHostTask(...) with HostTask has property: 1 -MHostTask->has_native_events(): 1 -Command::addDep -Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 -Command::addDep -Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 -Command::addDep -Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 -In CodeplayHostTask /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp2527 : 2527 -setHasNativeEvents(1); -Events.push_back: /home/hugh/llvm/sycl/source/detail/queue_impl.cpp212 -Command::addDep -Events.push_back: /home/hugh/llvm/sycl/source/detail/scheduler/commands.cpp642 -Making interop_handle with dep events: 0 -ih.get_native_events(): MEventImplPtrs.size() = 0 -native_events.size(): 0 -Have 0 native events in host_task 2! -Got native events from kernel: 1 -Getting native event dependencies for new host_task! -In setNativeEvents with native_events.size: 1 -Command has deps: 2 -success