From 023934fb8080698d1ddb4caa272fb58bfdb0524f Mon Sep 17 00:00:00 2001 From: Ben Tracy Date: Tue, 25 Jul 2023 18:35:34 +0100 Subject: [PATCH 01/10] [SYCL][Graph] Error when immediate command lists are used (#277) - Error when we detected immediate command lists - Throws exception with sycl::invalid - Test which uses property::queue::immediate_command_list to test errors. --- .../ur/adapters/level_zero/command_buffer.cpp | 6 +++ sycl/source/detail/graph_impl.cpp | 9 +++- .../Graph/immediate_command_list_error.cpp | 47 +++++++++++++++++++ 3 files changed, 61 insertions(+), 1 deletion(-) create mode 100644 sycl/test-e2e/Graph/immediate_command_list_error.cpp diff --git a/sycl/plugins/unified_runtime/ur/adapters/level_zero/command_buffer.cpp b/sycl/plugins/unified_runtime/ur/adapters/level_zero/command_buffer.cpp index edf1f7c81663f..509cc5c4ff418 100644 --- a/sycl/plugins/unified_runtime/ur/adapters/level_zero/command_buffer.cpp +++ b/sycl/plugins/unified_runtime/ur/adapters/level_zero/command_buffer.cpp @@ -661,6 +661,12 @@ UR_APIEXPORT ur_result_t UR_APICALL urCommandBufferEnqueueExp( ur_exp_command_buffer_handle_t CommandBuffer, ur_queue_handle_t Queue, uint32_t NumEventsInWaitList, const ur_event_handle_t *EventWaitList, ur_event_handle_t *Event) { + // There are issues with immediate command lists so return an error if the + // queue is in that mode. + if (Queue->UsingImmCmdLists) { + return UR_RESULT_ERROR_INVALID_QUEUE_PROPERTIES; + } + std::scoped_lock lock(Queue->Mutex); // Use compute engine rather than copy engine const auto UseCopyEngine = false; diff --git a/sycl/source/detail/graph_impl.cpp b/sycl/source/detail/graph_impl.cpp index 1a4f07285fe4b..26e34cbe7b15c 100644 --- a/sycl/source/detail/graph_impl.cpp +++ b/sycl/source/detail/graph_impl.cpp @@ -393,7 +393,14 @@ exec_graph_impl::enqueue(const std::shared_ptr &Queue, ->call_nocheck< sycl::detail::PiApiKind::piextEnqueueCommandBuffer>( CommandBuffer, Queue->getHandleRef(), 0, nullptr, OutEvent); - if (Res != pi_result::PI_SUCCESS) { + if (Res == pi_result::PI_ERROR_INVALID_QUEUE_PROPERTIES) { + throw sycl::exception( + make_error_code(errc::invalid), + "Graphs cannot be submitted to a queue which uses " + "immediate command lists. Use " + "sycl::ext::intel::property::queue::no_immediate_" + "command_list to disable them."); + } else if (Res != pi_result::PI_SUCCESS) { throw sycl::exception( errc::event, "Failed to enqueue event for command buffer submission"); diff --git a/sycl/test-e2e/Graph/immediate_command_list_error.cpp b/sycl/test-e2e/Graph/immediate_command_list_error.cpp new file mode 100644 index 0000000000000..bad3fac48007c --- /dev/null +++ b/sycl/test-e2e/Graph/immediate_command_list_error.cpp @@ -0,0 +1,47 @@ +// REQUIRES: level_zero, gpu +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out +// RUN: %if ext_oneapi_level_zero %{env ZE_DEBUG=4 %{run} %t.out 2>&1 | FileCheck %s %} +// +// CHECK-NOT: LEAK + +// Tests that graph submission will throw if the target queue is using immediate +// command lists and not throw if they are using regular command queues. + +#include "graph_common.hpp" + +int main() { + queue QueueImmediate{ + {sycl::ext::intel::property::queue::immediate_command_list{}}}; + queue QueueNoImmediate{ + QueueImmediate.get_context(), + QueueImmediate.get_device(), + {sycl::ext::intel::property::queue::no_immediate_command_list{}}}; + + exp_ext::command_graph Graph{QueueNoImmediate.get_context(), + QueueNoImmediate.get_device()}; + + std::error_code ErrorCode = make_error_code(sycl::errc::success); + try { + auto GraphExec = Graph.finalize(); + QueueNoImmediate.submit( + [&](handler &CGH) { CGH.ext_oneapi_graph(GraphExec); }); + } catch (sycl::exception &E) { + ErrorCode = E.code(); + } + + assert(ErrorCode == make_error_code(errc::success)); + + ErrorCode = make_error_code(sycl::errc::success); + try { + auto GraphExec = Graph.finalize(); + QueueImmediate.submit( + [&](handler &CGH) { CGH.ext_oneapi_graph(GraphExec); }); + } catch (sycl::exception &E) { + ErrorCode = E.code(); + } + + assert(ErrorCode == make_error_code(errc::invalid)); + + return 0; +} From 4403358e39567c74346dc78023b6b1bab965a68f Mon Sep 17 00:00:00 2001 From: Maxime France-Pillois Date: Thu, 3 Aug 2023 15:32:19 +0100 Subject: [PATCH 02/10] [SYCL][Graph] Implement exceptions for incompatible extensions (#276) * [SYCL][Graph] Implement exceptions for incompatible extensions Throws an invalid exception when trying to use the following extensions along with Graph. - sycl_ext_oneapi_enqueue_barrier - sycl_ext_oneapi_memcpy2d - sycl_ext_codeplay_kernel_fusion - sycl_ext_oneapi_kernel_properties - sycl_ext_oneapi_device_global Closes Issue: #154 * [SYCL][Graph] Implement exceptions for incompatible extensions Adds info to exception message Moves tests from e2e to unitests when possible * [SYCL][Graph] Implement exceptions for incompatible extensions Corrects some typos and adds comments. * [SYCL][Graph] Implement exceptions for incompatible extensions Used a template function to throw exception instead of a parametrized function. * [SYCL][Graph] Implement exceptions for incompatible extensions Moves Sycl-extension enum definition. Limits graph recording to non-explicit path in the new tests. * [SYCL][Graph] Implement exceptions for incompatible extensions Updates Linux ABI dump file with the new handler function throwing exception. --- sycl/include/sycl/handler.hpp | 44 ++- .../detail/fusion/fusion_wrapper_impl.cpp | 5 + sycl/source/detail/graph_impl.cpp | 6 + sycl/source/handler.cpp | 49 ++- ...raph_exception_global_device_extension.cpp | 150 +++++++ sycl/test/abi/sycl_symbols_linux.dump | 4 + sycl/unittests/Extensions/CommandGraph.cpp | 367 ++++++++++++++++++ 7 files changed, 623 insertions(+), 2 deletions(-) create mode 100644 sycl/test-e2e/Graph/graph_exception_global_device_extension.cpp diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 33cfcdc70057d..894ccffcf2467 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -134,8 +134,19 @@ class pipe; } namespace ext::oneapi::experimental::detail { +// List of sycl experimental extensions +// This enum is used to define the extension from which a function is called. +// This is used in handler::throwIfGraphAssociated() to specify +// the message of the thrown expection. +enum SyclExtensions { + sycl_ext_oneapi_kernel_properties, + sycl_ext_oneapi_enqueue_barrier, + sycl_ext_oneapi_memcpy2d, + sycl_ext_oneapi_device_global +}; + class graph_impl; -} +} // namespace ext::oneapi::experimental::detail namespace detail { class handler_impl; @@ -2107,6 +2118,7 @@ class __SYCL_EXPORT handler { std::enable_if_t< ext::oneapi::experimental::is_property_list::value> single_task(PropertiesT Props, _KERNELFUNCPARAM(KernelFunc)) { + throwIfGraphAssociatedAndKernelProperties(); single_task_lambda_impl(Props, KernelFunc); } @@ -2117,6 +2129,7 @@ class __SYCL_EXPORT handler { ext::oneapi::experimental::is_property_list::value> parallel_for(range<1> NumWorkItems, PropertiesT Props, _KERNELFUNCPARAM(KernelFunc)) { + throwIfGraphAssociatedAndKernelProperties(); parallel_for_lambda_impl( NumWorkItems, Props, std::move(KernelFunc)); } @@ -2127,6 +2140,7 @@ class __SYCL_EXPORT handler { ext::oneapi::experimental::is_property_list::value> parallel_for(range<2> NumWorkItems, PropertiesT Props, _KERNELFUNCPARAM(KernelFunc)) { + throwIfGraphAssociatedAndKernelProperties(); parallel_for_lambda_impl( NumWorkItems, Props, std::move(KernelFunc)); } @@ -2137,6 +2151,7 @@ class __SYCL_EXPORT handler { ext::oneapi::experimental::is_property_list::value> parallel_for(range<3> NumWorkItems, PropertiesT Props, _KERNELFUNCPARAM(KernelFunc)) { + throwIfGraphAssociatedAndKernelProperties(); parallel_for_lambda_impl( NumWorkItems, Props, std::move(KernelFunc)); } @@ -2147,6 +2162,7 @@ class __SYCL_EXPORT handler { ext::oneapi::experimental::is_property_list::value> parallel_for(nd_range Range, PropertiesT Properties, _KERNELFUNCPARAM(KernelFunc)) { + throwIfGraphAssociatedAndKernelProperties(); parallel_for_impl(Range, Properties, std::move(KernelFunc)); } @@ -2159,6 +2175,7 @@ class __SYCL_EXPORT handler { detail::AreAllButLastReductions::value && ext::oneapi::experimental::is_property_list::value> parallel_for(range<1> Range, PropertiesT Properties, RestT &&...Rest) { + throwIfGraphAssociatedAndKernelProperties(); detail::reduction_parallel_for(*this, Range, Properties, std::forward(Rest)...); } @@ -2170,6 +2187,7 @@ class __SYCL_EXPORT handler { detail::AreAllButLastReductions::value && ext::oneapi::experimental::is_property_list::value> parallel_for(range<2> Range, PropertiesT Properties, RestT &&...Rest) { + throwIfGraphAssociatedAndKernelProperties(); detail::reduction_parallel_for(*this, Range, Properties, std::forward(Rest)...); } @@ -2181,6 +2199,7 @@ class __SYCL_EXPORT handler { detail::AreAllButLastReductions::value && ext::oneapi::experimental::is_property_list::value> parallel_for(range<3> Range, PropertiesT Properties, RestT &&...Rest) { + throwIfGraphAssociatedAndKernelProperties(); detail::reduction_parallel_for(*this, Range, Properties, std::forward(Rest)...); } @@ -2235,6 +2254,7 @@ class __SYCL_EXPORT handler { int Dims, typename PropertiesT> void parallel_for_work_group(range NumWorkGroups, PropertiesT Props, _KERNELFUNCPARAM(KernelFunc)) { + throwIfGraphAssociatedAndKernelProperties(); parallel_for_work_group_lambda_impl(NumWorkGroups, Props, KernelFunc); @@ -2245,6 +2265,7 @@ class __SYCL_EXPORT handler { void parallel_for_work_group(range NumWorkGroups, range WorkGroupSize, PropertiesT Props, _KERNELFUNCPARAM(KernelFunc)) { + throwIfGraphAssociatedAndKernelProperties(); parallel_for_work_group_lambda_impl( NumWorkGroups, WorkGroupSize, Props, KernelFunc); @@ -2552,6 +2573,8 @@ class __SYCL_EXPORT handler { /// until all commands previously submitted to this queue have entered the /// complete state. void ext_oneapi_barrier() { + throwIfGraphAssociated(); throwIfActionIsCreated(); setType(detail::CG::Barrier); } @@ -2637,6 +2660,8 @@ class __SYCL_EXPORT handler { typename = std::enable_if_t>> void ext_oneapi_memcpy2d(void *Dest, size_t DestPitch, const void *Src, size_t SrcPitch, size_t Width, size_t Height) { + throwIfGraphAssociated(); throwIfActionIsCreated(); if (Width > DestPitch) throw sycl::exception(sycl::make_error_code(errc::invalid), @@ -2815,6 +2840,8 @@ class __SYCL_EXPORT handler { void memcpy(ext::oneapi::experimental::device_global &Dest, const void *Src, size_t NumBytes = sizeof(T), size_t DestOffset = 0) { + throwIfGraphAssociated(); if (sizeof(T) < DestOffset + NumBytes) throw sycl::exception(make_error_code(errc::invalid), "Copy to device_global is out of bounds."); @@ -2847,6 +2874,8 @@ class __SYCL_EXPORT handler { memcpy(void *Dest, const ext::oneapi::experimental::device_global &Src, size_t NumBytes = sizeof(T), size_t SrcOffset = 0) { + throwIfGraphAssociated(); if (sizeof(T) < SrcOffset + NumBytes) throw sycl::exception(make_error_code(errc::invalid), "Copy from device_global is out of bounds."); @@ -3368,8 +3397,21 @@ class __SYCL_EXPORT handler { "handler::require() before it can be used."); } + template + std::enable_if_t< + ext::oneapi::experimental::is_property_list::value> + throwIfGraphAssociatedAndKernelProperties() { + if (!std::is_same_v) + throwIfGraphAssociated(); + } + // Set value of the gpu cache configuration for the kernel. void setKernelCacheConfig(sycl::detail::pi::PiKernelCacheConfig); + + template + void throwIfGraphAssociated(); }; } // namespace _V1 } // namespace sycl diff --git a/sycl/source/detail/fusion/fusion_wrapper_impl.cpp b/sycl/source/detail/fusion/fusion_wrapper_impl.cpp index 492b0bc4aa852..d846b018ab64c 100644 --- a/sycl/source/detail/fusion/fusion_wrapper_impl.cpp +++ b/sycl/source/detail/fusion/fusion_wrapper_impl.cpp @@ -27,6 +27,11 @@ bool fusion_wrapper_impl::is_in_fusion_mode() const { } void fusion_wrapper_impl::start_fusion() { + if (MQueue->getCommandGraph()) { + throw sycl::exception(sycl::make_error_code(errc::invalid), + "SYCL kernel fusion can NOT be started " + "on a queue that is in a recording state."); + } detail::Scheduler::getInstance().startFusion(MQueue); } diff --git a/sycl/source/detail/graph_impl.cpp b/sycl/source/detail/graph_impl.cpp index 26e34cbe7b15c..143172c3ea80a 100644 --- a/sycl/source/detail/graph_impl.cpp +++ b/sycl/source/detail/graph_impl.cpp @@ -515,6 +515,12 @@ modifiable_command_graph::finalize(const sycl::property_list &) const { bool modifiable_command_graph::begin_recording(queue &RecordingQueue) { auto QueueImpl = sycl::detail::getSyclObjImpl(RecordingQueue); + if (QueueImpl->is_in_fusion_mode()) { + throw sycl::exception(sycl::make_error_code(errc::invalid), + "SYCL queue in kernel in fusion mode " + "can NOT be recorded."); + } + if (QueueImpl->get_context() != impl->getContext()) { throw sycl::exception(sycl::make_error_code(errc::invalid), "begin_recording called for a queue whose context " diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index 95db5c5eb66af..9ae123cf96379 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -794,6 +794,8 @@ void handler::verifyUsedKernelBundle(const std::string &KernelName) { } void handler::ext_oneapi_barrier(const std::vector &WaitList) { + throwIfGraphAssociated(); throwIfActionIsCreated(); MCGType = detail::CG::BarrierWaitlist; MEventsWaitWithBarrier.resize(WaitList.size()); @@ -1338,5 +1340,50 @@ handler::getCommandGraph() const { return MQueue->getCommandGraph(); } -} // namespace _V1 +template void handler::throwIfGraphAssociated< + ext::oneapi::experimental::detail::SyclExtensions:: + sycl_ext_oneapi_kernel_properties>(); +template void handler::throwIfGraphAssociated< + ext::oneapi::experimental::detail::SyclExtensions:: + sycl_ext_oneapi_enqueue_barrier>(); +template void +handler::throwIfGraphAssociated(); +template void handler::throwIfGraphAssociated< + ext::oneapi::experimental::detail::SyclExtensions:: + sycl_ext_oneapi_device_global>(); + +template +void handler::throwIfGraphAssociated() { + std::string ExceptionMsg = ""; + + if constexpr (ExtensionT == + ext::oneapi::experimental::detail::SyclExtensions:: + sycl_ext_oneapi_kernel_properties) { + ExceptionMsg = "sycl_ext_oneapi_kernel_properties"; + } + if constexpr (ExtensionT == + ext::oneapi::experimental::detail::SyclExtensions:: + sycl_ext_oneapi_enqueue_barrier) { + ExceptionMsg = "sycl_ext_oneapi_enqueue_barrier"; + } + if constexpr (ExtensionT == ext::oneapi::experimental::detail:: + SyclExtensions::sycl_ext_oneapi_memcpy2d) { + ExceptionMsg = "sycl_ext_oneapi_memcpy2d"; + } + if constexpr (ExtensionT == + ext::oneapi::experimental::detail::SyclExtensions:: + sycl_ext_oneapi_device_global) { + ExceptionMsg = "sycl_ext_oneapi_device_global"; + } + + if (MGraph || MQueue->getCommandGraph()) { + throw sycl::exception(sycl::make_error_code(errc::invalid), + "The feature " + ExceptionMsg + + " is not yet available " + "along with SYCL Graph extension."); + } +} + +} // __SYCL_INLINE_VER_NAMESPACE(_V1) } // namespace sycl diff --git a/sycl/test-e2e/Graph/graph_exception_global_device_extension.cpp b/sycl/test-e2e/Graph/graph_exception_global_device_extension.cpp new file mode 100644 index 0000000000000..e674beec4693f --- /dev/null +++ b/sycl/test-e2e/Graph/graph_exception_global_device_extension.cpp @@ -0,0 +1,150 @@ +// REQUIRES: level_zero, gpu +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out +// +// The test checks that invalid exception is thrown +// when trying to use sycl_ext_oneapi_device_global +// along with Graph. + +#include "graph_common.hpp" + +using TestProperties = decltype(sycl::ext::oneapi::experimental::properties{}); + +sycl::ext::oneapi::experimental::device_global + MemcpyDeviceGlobal; +sycl::ext::oneapi::experimental::device_global + CopyDeviceGlobal; + +enum OperationPath { Explicit, RecordReplay, Shortcut }; + +template void test() { + queue Q; + int MemcpyWrite = 42, CopyWrite = 24, MemcpyRead = 1, CopyRead = 2; + + exp_ext::command_graph Graph{Q.get_context(), Q.get_device()}; + + if constexpr (PathKind != OperationPath::Explicit) { + Graph.begin_recording(Q); + } + + // Copy from device globals before having written anything. + std::error_code ExceptionCode = make_error_code(sycl::errc::success); + try { + if constexpr (PathKind == OperationPath::Shortcut) { + Q.memcpy(&MemcpyRead, MemcpyDeviceGlobal); + } + if constexpr (PathKind == OperationPath::RecordReplay) { + Q.submit([&](handler &CGH) { + return CGH.memcpy(&MemcpyRead, MemcpyDeviceGlobal); + }); + } + if constexpr (PathKind == OperationPath::Explicit) { + Graph.add([&](handler &CGH) { + return CGH.memcpy(&MemcpyRead, MemcpyDeviceGlobal); + }); + } + } catch (exception &Exception) { + ExceptionCode = Exception.code(); + } + assert(ExceptionCode == sycl::errc::invalid); + + ExceptionCode = make_error_code(sycl::errc::success); + try { + if constexpr (PathKind == OperationPath::Shortcut) { + Q.copy(CopyDeviceGlobal, &CopyRead); + } + if constexpr (PathKind == OperationPath::RecordReplay) { + Q.submit( + [&](handler &CGH) { return CGH.copy(CopyDeviceGlobal, &CopyRead); }); + } + if constexpr (PathKind == OperationPath::Explicit) { + Graph.add( + [&](handler &CGH) { return CGH.copy(CopyDeviceGlobal, &CopyRead); }); + } + } catch (exception &Exception) { + ExceptionCode = Exception.code(); + } + assert(ExceptionCode == sycl::errc::invalid); + + // Write to device globals and then read their values. + ExceptionCode = make_error_code(sycl::errc::success); + try { + if constexpr (PathKind == OperationPath::Shortcut) { + Q.memcpy(MemcpyDeviceGlobal, &MemcpyWrite); + } + if constexpr (PathKind == OperationPath::RecordReplay) { + Q.submit([&](handler &CGH) { + return CGH.memcpy(MemcpyDeviceGlobal, &MemcpyWrite); + }); + } + if constexpr (PathKind == OperationPath::Explicit) { + Graph.add([&](handler &CGH) { + return CGH.memcpy(MemcpyDeviceGlobal, &MemcpyWrite); + }); + } + } catch (exception &Exception) { + ExceptionCode = Exception.code(); + } + assert(ExceptionCode == sycl::errc::invalid); + + ExceptionCode = make_error_code(sycl::errc::success); + try { + if constexpr (PathKind == OperationPath::Shortcut) { + Q.copy(&CopyWrite, CopyDeviceGlobal); + } else if constexpr (PathKind == OperationPath::RecordReplay) { + Q.submit( + [&](handler &CGH) { return CGH.copy(&CopyWrite, CopyDeviceGlobal); }); + } else if constexpr (PathKind == OperationPath::Explicit) { + Graph.add( + [&](handler &CGH) { return CGH.copy(&CopyWrite, CopyDeviceGlobal); }); + } + } catch (exception &Exception) { + ExceptionCode = Exception.code(); + } + assert(ExceptionCode == sycl::errc::invalid); + + ExceptionCode = make_error_code(sycl::errc::success); + try { + if constexpr (PathKind == OperationPath::Shortcut) { + Q.memcpy(&MemcpyRead, MemcpyDeviceGlobal); + } else if constexpr (PathKind == OperationPath::RecordReplay) { + Q.submit([&](handler &CGH) { + return CGH.memcpy(&MemcpyRead, MemcpyDeviceGlobal); + }); + } else if constexpr (PathKind == OperationPath::Explicit) { + Graph.add([&](handler &CGH) { + return CGH.memcpy(&MemcpyRead, MemcpyDeviceGlobal); + }); + } + } catch (exception &Exception) { + ExceptionCode = Exception.code(); + } + assert(ExceptionCode == sycl::errc::invalid); + + ExceptionCode = make_error_code(sycl::errc::success); + try { + if constexpr (PathKind == OperationPath::Shortcut) { + Q.copy(CopyDeviceGlobal, &CopyRead); + } else if constexpr (PathKind == OperationPath::RecordReplay) { + Q.submit( + [&](handler &CGH) { return CGH.copy(CopyDeviceGlobal, &CopyRead); }); + } else if constexpr (PathKind == OperationPath::Explicit) { + Graph.add( + [&](handler &CGH) { return CGH.copy(CopyDeviceGlobal, &CopyRead); }); + } + } catch (exception &Exception) { + ExceptionCode = Exception.code(); + } + assert(ExceptionCode == sycl::errc::invalid); + + if constexpr (PathKind != OperationPath::Explicit) { + Graph.end_recording(); + } +} + +int main() { + test(); + test(); + test(); + return 0; +} diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 9275be530451f..42373dc301c24 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -4098,6 +4098,10 @@ _ZN4sycl3_V17handler22memcpyFromDeviceGlobalEPvPKvbmm _ZN4sycl3_V17handler22setHandlerKernelBundleENS0_6kernelE _ZN4sycl3_V17handler22setHandlerKernelBundleERKSt10shared_ptrINS0_6detail18kernel_bundle_implEE _ZN4sycl3_V17handler22verifyUsedKernelBundleERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE +_ZN4sycl3_V17handler22throwIfGraphAssociatedILNS0_3ext6oneapi12experimental6detail14SyclExtensionsE0EEEvv +_ZN4sycl3_V17handler22throwIfGraphAssociatedILNS0_3ext6oneapi12experimental6detail14SyclExtensionsE1EEEvv +_ZN4sycl3_V17handler22throwIfGraphAssociatedILNS0_3ext6oneapi12experimental6detail14SyclExtensionsE3EEEvv +_ZN4sycl3_V17handler22throwIfGraphAssociatedILNS0_3ext6oneapi12experimental6detail14SyclExtensionsE2EEEvv _ZN4sycl3_V17handler24GetRangeRoundingSettingsERmS2_S2_ _ZN4sycl3_V17handler24ext_intel_read_host_pipeERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPvmb _ZN4sycl3_V17handler24ext_oneapi_memcpy2d_implEPvmPKvmmm diff --git a/sycl/unittests/Extensions/CommandGraph.cpp b/sycl/unittests/Extensions/CommandGraph.cpp index 590d86b8e0019..10cac8b18dcbd 100644 --- a/sycl/unittests/Extensions/CommandGraph.cpp +++ b/sycl/unittests/Extensions/CommandGraph.cpp @@ -20,6 +20,279 @@ using namespace sycl; using namespace sycl::ext::oneapi; +namespace { +/// Define the three possible path to add node to a SYCL Graph. +/// Shortcut is a sub-type of Record&Replay using Queue shortcut +/// instead of standard kernel submitions. +enum OperationPath { Explicit, RecordReplay, Shortcut }; + +/// Function types and classes for testing Kernel with properties extension +enum class Variant { Function, Functor, FunctorAndProperty }; + +template +class ReqdWGSizePositiveA; +template class ReqPositiveA; + +template range repeatRange(size_t Val); +template <> range<1> repeatRange<1>(size_t Val) { return range<1>{Val}; } +template <> range<2> repeatRange<2>(size_t Val) { return range<2>{Val, Val}; } +template <> range<3> repeatRange<3>(size_t Val) { + return range<3>{Val, Val, Val}; +} + +template struct KernelFunctorWithWGSizeProp { + void operator()(nd_item) const {} + void operator()(item) const {} + + auto get(sycl::ext::oneapi::experimental::properties_tag) { + return sycl::ext::oneapi::experimental::properties{ + sycl::ext::oneapi::experimental::work_group_size}; + } +}; + +/// Tries to add a Parallel_for node with kernel properties to the graph G +/// It tests that an invalid exception has been thrown +/// Since sycl_ext_oneapi_kernel_properties extension can not be used +/// along with SYCL Graph. +/// +/// @param G Modifiable graph to add commands to. +/// @param Q Queue to submit nodes to. +/// @param Props Properties associated to the submitted kernel +/// @param KernelFunc pointer to the kernel +template +void addKernelWithProperties( + sycl::ext::oneapi::experimental::detail::modifiable_command_graph &G, + queue &Q, PropertiesT Props, KernelType KernelFunc) { + constexpr size_t Dims = sizeof...(Is); + + // Test Parallel_for + std::error_code ExceptionCode = make_error_code(sycl::errc::success); + try { + if constexpr (PathKind == OperationPath::RecordReplay) { + Q.submit([&](handler &CGH) { + CGH.parallel_for>( + nd_range(repeatRange(8), range(Is...)), Props, + KernelFunc); + }); + } + if constexpr (PathKind == OperationPath::Shortcut) { + Q.parallel_for>( + nd_range(repeatRange(8), range(Is...)), Props, + KernelFunc); + } + if constexpr (PathKind == OperationPath::Explicit) { + G.add([&](handler &CGH) { + CGH.parallel_for>( + nd_range(repeatRange(8), range(Is...)), Props, + KernelFunc); + }); + } + } catch (exception &Exception) { + ExceptionCode = Exception.code(); + } + ASSERT_EQ(ExceptionCode, sycl::errc::invalid); +} + +/// Tries to add a Single task node with kernel properties to the graph G +/// It tests that an invalid exception has been thrown +/// Since sycl_ext_oneapi_kernel_properties extension can not be used +/// along with SYCL Graph. +/// +/// @param G Modifiable graph to add commands to. +/// @param Q Queue to submit nodes to. +/// @param Props Properties associated to the submitted kernel +/// @param KernelFunc pointer to the kernel +template +void testSingleTaskProperties(experimental::detail::modifiable_command_graph &G, + queue &Q, PropertiesT Props, + KernelType KernelFunc) { + + // Test Single_task + std::error_code ExceptionCode = make_error_code(sycl::errc::success); + try { + if constexpr (PathKind == OperationPath::RecordReplay) { + G.begin_recording(Q); + Q.submit([&](sycl::handler &CGH) { + CGH.single_task>(Props, + KernelFunc); + }); + G.end_recording(); + } + if constexpr (PathKind == OperationPath::Explicit) { + G.add([&](sycl::handler &CGH) { + CGH.single_task>(Props, + KernelFunc); + }); + } + } catch (exception &Exception) { + ExceptionCode = Exception.code(); + } + ASSERT_EQ(ExceptionCode, sycl::errc::invalid); +} + +/// This function groups all the different test cases +/// when adding a Parallel_for node with kernel properties to the graph G +/// +/// @param G Modifiable graph to add commands to. +/// @param Q Queue to submit nodes to. +template +void testParallelForProperties( + queue &Q, experimental::detail::modifiable_command_graph &G) { + auto Props = ext::oneapi::experimental::properties{ + experimental::work_group_size}; + auto KernelFunction = [](auto) {}; + + KernelFunctorWithWGSizeProp KernelFunctor; + + G.begin_recording(Q); + + addKernelWithProperties(G, Q, Props, KernelFunction); + addKernelWithProperties(G, Q, Props, + KernelFunctor); + + addKernelWithProperties( + G, Q, Props, KernelFunction); + addKernelWithProperties(G, Q, Props, KernelFunctor); + + G.end_recording(); + + addKernelWithProperties( + G, Q, Props, KernelFunction); + addKernelWithProperties(G, Q, Props, KernelFunctor); +} + +/// Tries to enqueue oneapi barrier to the graph G +/// It tests that an invalid exception has been thrown +/// Since sycl_ext_oneapi_enqueue_barrier extension can not be used +/// along with SYCL Graph. +template void testEnqueueBarrier() { + sycl::context Context; + sycl::queue Q1(Context, sycl::default_selector_v); + + experimental::command_graph Graph1{ + Q1.get_context(), Q1.get_device()}; + + Graph1.add([&](sycl::handler &cgh) {}); + Graph1.add([&](sycl::handler &cgh) {}); + + if constexpr (PathKind != OperationPath::Explicit) { + Graph1.begin_recording(Q1); + } + + // call queue::ext_oneapi_submit_barrier() + std::error_code ExceptionCode = make_error_code(sycl::errc::success); + try { + if constexpr (PathKind == OperationPath::Shortcut) { + Q1.ext_oneapi_submit_barrier(); + } + if constexpr (PathKind == OperationPath::RecordReplay) { + Q1.submit([&](sycl::handler &CGH) { CGH.ext_oneapi_barrier(); }); + } + if constexpr (PathKind == OperationPath::Explicit) { + Graph1.add([&](handler &CGH) { CGH.ext_oneapi_barrier(); }); + } + + } catch (exception &Exception) { + ExceptionCode = Exception.code(); + } + ASSERT_EQ(ExceptionCode, sycl::errc::invalid); + + if constexpr (PathKind != OperationPath::Explicit) { + Graph1.end_recording(); + } + + sycl::queue Q2(Context, sycl::default_selector_v); + sycl::queue Q3(Context, sycl::default_selector_v); + + experimental::command_graph Graph2{ + Q2.get_context(), Q2.get_device()}; + experimental::command_graph Graph3{ + Q3.get_context(), Q3.get_device()}; + + Graph2.begin_recording(Q2); + Graph3.begin_recording(Q3); + + auto Event1 = Q2.submit([&](sycl::handler &cgh) {}); + auto Event2 = Q3.submit([&](sycl::handler &cgh) {}); + + if constexpr (PathKind == OperationPath::Explicit) { + Graph2.end_recording(); + Graph3.end_recording(); + } + + // call handler::barrier(const std::vector &WaitList) + ExceptionCode = make_error_code(sycl::errc::success); + try { + if constexpr (PathKind == OperationPath::Shortcut) { + Q3.ext_oneapi_submit_barrier({Event1, Event2}); + } + if constexpr (PathKind == OperationPath::RecordReplay) { + Q3.submit([&](sycl::handler &CGH) { + CGH.ext_oneapi_barrier({Event1, Event2}); + }); + } + if constexpr (PathKind == OperationPath::Explicit) { + Graph3.add([&](handler &CGH) { + CGH.ext_oneapi_barrier({Event1, Event2}); + }); + } + + } catch (exception &Exception) { + ExceptionCode = Exception.code(); + } + ASSERT_EQ(ExceptionCode, sycl::errc::invalid); + + if constexpr (PathKind != OperationPath::Explicit) { + Graph2.end_recording(); + Graph3.end_recording(); + } +} + +/// Tries to add a memcpy2D node to the graph G +/// It tests that an invalid exception has been thrown +/// Since sycl_ext_oneapi_memcpy2d extension can not be used +/// along with SYCL Graph. +/// +/// @param G Modifiable graph to add commands to. +/// @param Q Queue to submit nodes to. +/// @param Dest Pointer to the memory destination +/// @param DestPitch pitch at the destination +/// @param Src Pointer to the memory source +/// @param SrcPitch pitch at the source +/// @param Witdh width of the data to copy +/// @param Height height of the data to copy +template +void addMemcpy2D(experimental::detail::modifiable_command_graph &G, queue &Q, + void *Dest, size_t DestPitch, const void *Src, size_t SrcPitch, + size_t Width, size_t Height) { + std::error_code ExceptionCode = make_error_code(sycl::errc::success); + try { + if constexpr (PathKind == OperationPath::RecordReplay) { + Q.submit([&](handler &CGH) { + CGH.ext_oneapi_memcpy2d(Dest, DestPitch, Src, SrcPitch, Width, Height); + }); + } + if constexpr (PathKind == OperationPath::Shortcut) { + Q.ext_oneapi_memcpy2d(Dest, DestPitch, Src, SrcPitch, Width, Height); + } + if constexpr (PathKind == OperationPath::Explicit) { + G.add([&](handler &CGH) { + CGH.ext_oneapi_memcpy2d(Dest, DestPitch, Src, SrcPitch, Width, Height); + }); + } + } catch (exception &Exception) { + ExceptionCode = Exception.code(); + } + ASSERT_EQ(ExceptionCode, sycl::errc::invalid); +} + +} // namespace + class CommandGraphTest : public ::testing::Test { public: CommandGraphTest() @@ -606,3 +879,97 @@ TEST_F(CommandGraphTest, InOrderQueueWithEmptyLast) { ASSERT_EQ(*ScheduleIt, PtrNode2); ASSERT_EQ(InOrderQueue.get_context(), GraphExecImpl->getContext()); } + +TEST_F(CommandGraphTest, EnqueueBarrierExceptionCheck) { + testEnqueueBarrier(); + testEnqueueBarrier(); + testEnqueueBarrier(); +} + +TEST_F(CommandGraphTest, FusionExtensionExceptionCheck) { + queue Q{ext::codeplay::experimental::property::queue::enable_fusion{}}; + + experimental::command_graph Graph{ + Q.get_context(), Q.get_device()}; + + ext::codeplay::experimental::fusion_wrapper fw{Q}; + + // Test: Start fusion on a queue that is in recording mode + Graph.begin_recording(Q); + + std::error_code ExceptionCode = make_error_code(sycl::errc::success); + try { + fw.start_fusion(); + } catch (exception &Exception) { + ExceptionCode = Exception.code(); + } + ASSERT_EQ(ExceptionCode, sycl::errc::invalid); + + Graph.end_recording(Q); + + // Test: begin recording a queue in fusion mode + + fw.start_fusion(); + + ExceptionCode = make_error_code(sycl::errc::success); + try { + Graph.begin_recording(Q); + } catch (exception &Exception) { + ExceptionCode = Exception.code(); + } + ASSERT_EQ(ExceptionCode, sycl::errc::invalid); +} + +TEST_F(CommandGraphTest, KernelPropertiesExceptionCheck) { + + // Test Parallel for entry point + testParallelForProperties<4>(Queue, Graph); + testParallelForProperties<4, 4>(Queue, Graph); + testParallelForProperties<8, 4>(Queue, Graph); + testParallelForProperties<4, 8>(Queue, Graph); + testParallelForProperties<4, 4, 4>(Queue, Graph); + testParallelForProperties<4, 4, 8>(Queue, Graph); + testParallelForProperties<8, 4, 4>(Queue, Graph); + testParallelForProperties<4, 8, 4>(Queue, Graph); + + // Test Single Task entry point + auto Props = ext::oneapi::experimental::properties{ + ext::oneapi::experimental::work_group_size<4>}; + auto KernelFunction = [](auto) {}; + testSingleTaskProperties(Graph, Queue, Props, + KernelFunction); + testSingleTaskProperties(Graph, Queue, Props, + KernelFunction); +} + +TEST_F(CommandGraphTest, Memcpy2DExceptionCheck) { + constexpr size_t RECT_WIDTH = 30; + constexpr size_t RECT_HEIGHT = 21; + constexpr size_t SRC_ELEMS = RECT_WIDTH * RECT_HEIGHT; + constexpr size_t DST_ELEMS = SRC_ELEMS; + + using T = int; + + Graph.begin_recording(Queue); + + T *USMMemSrc = malloc_device(SRC_ELEMS, Queue); + T *USMMemDst = malloc_device(DST_ELEMS, Queue); + + addMemcpy2D( + Graph, Queue, USMMemDst, RECT_WIDTH * sizeof(T), USMMemSrc, + RECT_WIDTH * sizeof(T), RECT_WIDTH * sizeof(T), RECT_HEIGHT); + + addMemcpy2D( + Graph, Queue, USMMemDst, RECT_WIDTH * sizeof(T), USMMemSrc, + RECT_WIDTH * sizeof(T), RECT_WIDTH * sizeof(T), RECT_HEIGHT); + + Graph.end_recording(); + + addMemcpy2D( + Graph, Queue, USMMemDst, RECT_WIDTH * sizeof(T), USMMemSrc, + RECT_WIDTH * sizeof(T), RECT_WIDTH * sizeof(T), RECT_HEIGHT); + + sycl::free(USMMemSrc, Queue); + sycl::free(USMMemDst, Queue); +} +>>>>>>> b99238b99781... [SYCL][Graph] Implement exceptions for incompatible extensions (#276) From 7d8e94a7cd6c9605bf411a895b96d21d993b342c Mon Sep 17 00:00:00 2001 From: Ben Tracy Date: Thu, 10 Aug 2023 15:34:40 +0100 Subject: [PATCH 03/10] [SYCL][Graph] Add exceptions on using spec constants, kernel bundles and reductions (#290) - Add exceptions when using spec constants, reductions and kernel bundles - Unit tests added for these. - Refactored handler throwing code to remove templates specializations. - Moved and renamed unsupported features enum to graph header. - Updated reduction tests to be XFAIL - Tweak exception message wording - Update ABI symbol tests. --------- Co-authored-by: Ewan Crawford --- .../sycl/ext/oneapi/experimental/graph.hpp | 35 ++++++++ sycl/include/sycl/handler.hpp | 70 +++++++++------ sycl/source/handler.cpp | 55 ++---------- .../RecordReplay/dotp_buffer_reduction.cpp | 3 + .../Graph/RecordReplay/dotp_usm_reduction.cpp | 3 + .../RecordReplay/sub_graph_reduction.cpp | 3 + sycl/test/abi/sycl_symbols_linux.dump | 6 +- sycl/test/abi/sycl_symbols_windows.dump | 2 +- sycl/unittests/Extensions/CommandGraph.cpp | 86 ++++++++++++++++++- 9 files changed, 184 insertions(+), 79 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/graph.hpp b/sycl/include/sycl/ext/oneapi/experimental/graph.hpp index 80837a48f07d7..3541cb62f1969 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/graph.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/graph.hpp @@ -31,6 +31,41 @@ namespace oneapi { namespace experimental { namespace detail { +// List of sycl features and extensions which are not supported by graphs. Used +// for throwing errors when these features are used with graphs. +enum class UnsupportedGraphFeatures { + sycl_reductions, + sycl_specialization_constants, + sycl_kernel_bundle, + sycl_ext_oneapi_kernel_properties, + sycl_ext_oneapi_enqueue_barrier, + sycl_ext_oneapi_memcpy2d, + sycl_ext_oneapi_device_global +}; + +constexpr const char * +UnsupportedFeatureToString(UnsupportedGraphFeatures Feature) { + using UGF = UnsupportedGraphFeatures; + switch (Feature) { + case UGF::sycl_reductions: + return "Reductions"; + case UGF::sycl_specialization_constants: + return "Specialization Constants"; + case UGF::sycl_kernel_bundle: + return "Kernel Bundles"; + case UGF::sycl_ext_oneapi_kernel_properties: + return "sycl_ext_oneapi_kernel_properties"; + case UGF::sycl_ext_oneapi_enqueue_barrier: + return "sycl_ext_oneapi_enqueue_barrier"; + case UGF::sycl_ext_oneapi_memcpy2d: + return "sycl_ext_oneapi_memcpy2d"; + case UGF::sycl_ext_oneapi_device_global: + return "sycl_ext_oneapi_device_global"; + default: + return {}; + } +} + class node_impl; class graph_impl; class exec_graph_impl; diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 894ccffcf2467..9d5a50c876ac3 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -134,17 +134,6 @@ class pipe; } namespace ext::oneapi::experimental::detail { -// List of sycl experimental extensions -// This enum is used to define the extension from which a function is called. -// This is used in handler::throwIfGraphAssociated() to specify -// the message of the thrown expection. -enum SyclExtensions { - sycl_ext_oneapi_kernel_properties, - sycl_ext_oneapi_enqueue_barrier, - sycl_ext_oneapi_memcpy2d, - sycl_ext_oneapi_device_global -}; - class graph_impl; } // namespace ext::oneapi::experimental::detail namespace detail { @@ -1589,6 +1578,10 @@ class __SYCL_EXPORT handler { void set_specialization_constant( typename std::remove_reference_t::value_type Value) { + throwIfGraphAssociated< + ext::oneapi::experimental::detail::UnsupportedGraphFeatures:: + sycl_specialization_constants>(); + setStateSpecConstSet(); std::shared_ptr KernelBundleImplPtr = @@ -1603,6 +1596,10 @@ class __SYCL_EXPORT handler { typename std::remove_reference_t::value_type get_specialization_constant() const { + throwIfGraphAssociated< + ext::oneapi::experimental::detail::UnsupportedGraphFeatures:: + sycl_specialization_constants>(); + if (isStateExplicitKernelBundle()) throw sycl::exception(make_error_code(errc::invalid), "Specialization constants cannot be read after " @@ -2175,6 +2172,8 @@ class __SYCL_EXPORT handler { detail::AreAllButLastReductions::value && ext::oneapi::experimental::is_property_list::value> parallel_for(range<1> Range, PropertiesT Properties, RestT &&...Rest) { + throwIfGraphAssociated(); throwIfGraphAssociatedAndKernelProperties(); detail::reduction_parallel_for(*this, Range, Properties, std::forward(Rest)...); @@ -2187,6 +2186,8 @@ class __SYCL_EXPORT handler { detail::AreAllButLastReductions::value && ext::oneapi::experimental::is_property_list::value> parallel_for(range<2> Range, PropertiesT Properties, RestT &&...Rest) { + throwIfGraphAssociated(); throwIfGraphAssociatedAndKernelProperties(); detail::reduction_parallel_for(*this, Range, Properties, std::forward(Rest)...); @@ -2199,6 +2200,8 @@ class __SYCL_EXPORT handler { detail::AreAllButLastReductions::value && ext::oneapi::experimental::is_property_list::value> parallel_for(range<3> Range, PropertiesT Properties, RestT &&...Rest) { + throwIfGraphAssociated(); throwIfGraphAssociatedAndKernelProperties(); detail::reduction_parallel_for(*this, Range, Properties, std::forward(Rest)...); @@ -2235,6 +2238,8 @@ class __SYCL_EXPORT handler { detail::AreAllButLastReductions::value && ext::oneapi::experimental::is_property_list::value> parallel_for(nd_range Range, PropertiesT Properties, RestT &&...Rest) { + throwIfGraphAssociated(); detail::reduction_parallel_for(*this, Range, Properties, std::forward(Rest)...); } @@ -2573,8 +2578,9 @@ class __SYCL_EXPORT handler { /// until all commands previously submitted to this queue have entered the /// complete state. void ext_oneapi_barrier() { - throwIfGraphAssociated(); + throwIfGraphAssociated< + ext::oneapi::experimental::detail::UnsupportedGraphFeatures:: + sycl_ext_oneapi_enqueue_barrier>(); throwIfActionIsCreated(); setType(detail::CG::Barrier); } @@ -2660,8 +2666,9 @@ class __SYCL_EXPORT handler { typename = std::enable_if_t>> void ext_oneapi_memcpy2d(void *Dest, size_t DestPitch, const void *Src, size_t SrcPitch, size_t Width, size_t Height) { - throwIfGraphAssociated(); + throwIfGraphAssociated< + ext::oneapi::experimental::detail::UnsupportedGraphFeatures:: + sycl_ext_oneapi_memcpy2d>(); throwIfActionIsCreated(); if (Width > DestPitch) throw sycl::exception(sycl::make_error_code(errc::invalid), @@ -2840,8 +2847,9 @@ class __SYCL_EXPORT handler { void memcpy(ext::oneapi::experimental::device_global &Dest, const void *Src, size_t NumBytes = sizeof(T), size_t DestOffset = 0) { - throwIfGraphAssociated(); + throwIfGraphAssociated< + ext::oneapi::experimental::detail::UnsupportedGraphFeatures:: + sycl_ext_oneapi_device_global>(); if (sizeof(T) < DestOffset + NumBytes) throw sycl::exception(make_error_code(errc::invalid), "Copy to device_global is out of bounds."); @@ -2874,8 +2882,9 @@ class __SYCL_EXPORT handler { memcpy(void *Dest, const ext::oneapi::experimental::device_global &Src, size_t NumBytes = sizeof(T), size_t SrcOffset = 0) { - throwIfGraphAssociated(); + throwIfGraphAssociated< + ext::oneapi::experimental::detail::UnsupportedGraphFeatures:: + sycl_ext_oneapi_device_global>(); if (sizeof(T) < SrcOffset + NumBytes) throw sycl::exception(make_error_code(errc::invalid), "Copy from device_global is out of bounds."); @@ -3400,18 +3409,31 @@ class __SYCL_EXPORT handler { template std::enable_if_t< ext::oneapi::experimental::is_property_list::value> - throwIfGraphAssociatedAndKernelProperties() { + throwIfGraphAssociatedAndKernelProperties() const { if (!std::is_same_v) - throwIfGraphAssociated(); + throwIfGraphAssociated< + ext::oneapi::experimental::detail::UnsupportedGraphFeatures:: + sycl_ext_oneapi_kernel_properties>(); } // Set value of the gpu cache configuration for the kernel. void setKernelCacheConfig(sycl::detail::pi::PiKernelCacheConfig); - template - void throwIfGraphAssociated(); + template < + ext::oneapi::experimental::detail::UnsupportedGraphFeatures FeatureT> + void throwIfGraphAssociated() const { + + if (getCommandGraph()) { + std::string FeatureString = + ext::oneapi::experimental::detail::UnsupportedFeatureToString( + FeatureT); + throw sycl::exception(sycl::make_error_code(errc::invalid), + "The " + FeatureString + + " feature is not yet available " + "for use with the SYCL Graph extension."); + } + } }; } // namespace _V1 } // namespace sycl diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index 9ae123cf96379..69036e2dfa190 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -794,8 +794,9 @@ void handler::verifyUsedKernelBundle(const std::string &KernelName) { } void handler::ext_oneapi_barrier(const std::vector &WaitList) { - throwIfGraphAssociated(); + throwIfGraphAssociated< + ext::oneapi::experimental::detail::UnsupportedGraphFeatures:: + sycl_ext_oneapi_enqueue_barrier>(); throwIfActionIsCreated(); MCGType = detail::CG::BarrierWaitlist; MEventsWaitWithBarrier.resize(WaitList.size()); @@ -1104,6 +1105,9 @@ void handler::ext_oneapi_signal_external_semaphore( void handler::use_kernel_bundle( const kernel_bundle &ExecBundle) { + throwIfGraphAssociated(); + std::shared_ptr PrimaryQueue = MImpl->MSubmissionPrimaryQueue; if (PrimaryQueue->get_context() != ExecBundle.get_context()) @@ -1340,50 +1344,5 @@ handler::getCommandGraph() const { return MQueue->getCommandGraph(); } -template void handler::throwIfGraphAssociated< - ext::oneapi::experimental::detail::SyclExtensions:: - sycl_ext_oneapi_kernel_properties>(); -template void handler::throwIfGraphAssociated< - ext::oneapi::experimental::detail::SyclExtensions:: - sycl_ext_oneapi_enqueue_barrier>(); -template void -handler::throwIfGraphAssociated(); -template void handler::throwIfGraphAssociated< - ext::oneapi::experimental::detail::SyclExtensions:: - sycl_ext_oneapi_device_global>(); - -template -void handler::throwIfGraphAssociated() { - std::string ExceptionMsg = ""; - - if constexpr (ExtensionT == - ext::oneapi::experimental::detail::SyclExtensions:: - sycl_ext_oneapi_kernel_properties) { - ExceptionMsg = "sycl_ext_oneapi_kernel_properties"; - } - if constexpr (ExtensionT == - ext::oneapi::experimental::detail::SyclExtensions:: - sycl_ext_oneapi_enqueue_barrier) { - ExceptionMsg = "sycl_ext_oneapi_enqueue_barrier"; - } - if constexpr (ExtensionT == ext::oneapi::experimental::detail:: - SyclExtensions::sycl_ext_oneapi_memcpy2d) { - ExceptionMsg = "sycl_ext_oneapi_memcpy2d"; - } - if constexpr (ExtensionT == - ext::oneapi::experimental::detail::SyclExtensions:: - sycl_ext_oneapi_device_global) { - ExceptionMsg = "sycl_ext_oneapi_device_global"; - } - - if (MGraph || MQueue->getCommandGraph()) { - throw sycl::exception(sycl::make_error_code(errc::invalid), - "The feature " + ExceptionMsg + - " is not yet available " - "along with SYCL Graph extension."); - } -} - -} // __SYCL_INLINE_VER_NAMESPACE(_V1) +} // namespace _V1 } // namespace sycl diff --git a/sycl/test-e2e/Graph/RecordReplay/dotp_buffer_reduction.cpp b/sycl/test-e2e/Graph/RecordReplay/dotp_buffer_reduction.cpp index af577686832cd..3e72076841306 100644 --- a/sycl/test-e2e/Graph/RecordReplay/dotp_buffer_reduction.cpp +++ b/sycl/test-e2e/Graph/RecordReplay/dotp_buffer_reduction.cpp @@ -6,6 +6,9 @@ // // CHECK-NOT: LEAK +// Expected fail as reduction support is not complete. +// XFAIL: * + #define GRAPH_E2E_RECORD_REPLAY #include "../Inputs/dotp_buffer_reduction.cpp" diff --git a/sycl/test-e2e/Graph/RecordReplay/dotp_usm_reduction.cpp b/sycl/test-e2e/Graph/RecordReplay/dotp_usm_reduction.cpp index dab4b34eec79d..6738affa87c13 100644 --- a/sycl/test-e2e/Graph/RecordReplay/dotp_usm_reduction.cpp +++ b/sycl/test-e2e/Graph/RecordReplay/dotp_usm_reduction.cpp @@ -6,6 +6,9 @@ // // CHECK-NOT: LEAK +// Expected fail as reduction support is not complete. +// XFAIL: * + #define GRAPH_E2E_RECORD_REPLAY #include "../Inputs/dotp_usm_reduction.cpp" diff --git a/sycl/test-e2e/Graph/RecordReplay/sub_graph_reduction.cpp b/sycl/test-e2e/Graph/RecordReplay/sub_graph_reduction.cpp index f3e58b1ef99ff..de9cbead9634d 100644 --- a/sycl/test-e2e/Graph/RecordReplay/sub_graph_reduction.cpp +++ b/sycl/test-e2e/Graph/RecordReplay/sub_graph_reduction.cpp @@ -6,6 +6,9 @@ // // CHECK-NOT: LEAK +// Expected fail as reduction support is not complete. +// XFAIL: * + #define GRAPH_E2E_RECORD_REPLAY #include "../Inputs/sub_graph_reduction.cpp" diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 42373dc301c24..1967793fea90a 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -4003,6 +4003,7 @@ _ZN4sycl3_V16detail30UnsampledImageAccessorBaseHost10getAccDataEv _ZN4sycl3_V16detail30UnsampledImageAccessorBaseHost6getPtrEv _ZN4sycl3_V16detail30UnsampledImageAccessorBaseHostC1ENS0_5rangeILi3EEENS0_6access4modeEPviiNS0_2idILi3EEENS0_18image_channel_typeENS0_19image_channel_orderERKNS0_13property_listE _ZN4sycl3_V16detail30UnsampledImageAccessorBaseHostC2ENS0_5rangeILi3EEENS0_6access4modeEPviiNS0_2idILi3EEENS0_18image_channel_typeENS0_19image_channel_orderERKNS0_13property_listE +_ZN4sycl3_V16detail33enable_ext_oneapi_default_contextEb _ZN4sycl3_V16detail33reduGetMaxNumConcurrentWorkGroupsESt10shared_ptrINS1_10queue_implEE _ZN4sycl3_V16detail34addHostSampledImageAccessorAndWaitEPNS1_28SampledImageAccessorImplHostE _ZN4sycl3_V16detail35sampledImageConstructorNotificationEPvS2_RKSt8optionalINS0_12image_targetEEPKvjRKNS1_13code_locationE @@ -4098,10 +4099,6 @@ _ZN4sycl3_V17handler22memcpyFromDeviceGlobalEPvPKvbmm _ZN4sycl3_V17handler22setHandlerKernelBundleENS0_6kernelE _ZN4sycl3_V17handler22setHandlerKernelBundleERKSt10shared_ptrINS0_6detail18kernel_bundle_implEE _ZN4sycl3_V17handler22verifyUsedKernelBundleERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE -_ZN4sycl3_V17handler22throwIfGraphAssociatedILNS0_3ext6oneapi12experimental6detail14SyclExtensionsE0EEEvv -_ZN4sycl3_V17handler22throwIfGraphAssociatedILNS0_3ext6oneapi12experimental6detail14SyclExtensionsE1EEEvv -_ZN4sycl3_V17handler22throwIfGraphAssociatedILNS0_3ext6oneapi12experimental6detail14SyclExtensionsE3EEEvv -_ZN4sycl3_V17handler22throwIfGraphAssociatedILNS0_3ext6oneapi12experimental6detail14SyclExtensionsE2EEEvv _ZN4sycl3_V17handler24GetRangeRoundingSettingsERmS2_S2_ _ZN4sycl3_V17handler24ext_intel_read_host_pipeERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEPvmb _ZN4sycl3_V17handler24ext_oneapi_memcpy2d_implEPvmPKvmmm @@ -4722,7 +4719,6 @@ _ZNK4sycl3_V18platform11get_backendEv _ZNK4sycl3_V18platform11get_devicesENS0_4info11device_typeE _ZNK4sycl3_V18platform13has_extensionERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE _ZNK4sycl3_V18platform30ext_oneapi_get_default_contextEv -_ZN4sycl3_V16detail33enable_ext_oneapi_default_contextEb _ZNK4sycl3_V18platform3getEv _ZNK4sycl3_V18platform3hasENS0_6aspectE _ZNK4sycl3_V18platform7is_hostEv diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index 545bdb4361454..d51c06d332cb0 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -979,6 +979,7 @@ ?die@pi@detail@_V1@sycl@@YAXPEBD@Z ?discard_or_return@queue@_V1@sycl@@AEAA?AVevent@23@AEBV423@@Z ?empty@kernel_bundle_plain@detail@_V1@sycl@@QEBA_NXZ +?enable_ext_oneapi_default_context@detail@_V1@sycl@@YAX_N@Z ?end@HostProfilingInfo@detail@_V1@sycl@@QEAAXXZ ?end@exception_list@_V1@sycl@@QEBA?AV?$_Vector_const_iterator@V?$_Vector_val@U?$_Simple_types@Vexception_ptr@std@@@std@@@std@@@std@@XZ ?end@kernel_bundle_plain@detail@_V1@sycl@@IEBAPEBVdevice_image_plain@234@XZ @@ -1025,7 +1026,6 @@ ?ext_oneapi_enable_peer_access@device@_V1@sycl@@QEAAXAEBV123@@Z ?ext_oneapi_fill2d_impl@handler@_V1@sycl@@AEAAXPEAX_KPEBX111@Z ?ext_oneapi_get_default_context@platform@_V1@sycl@@QEBA?AVcontext@23@XZ -?enable_ext_oneapi_default_context@detail@_V1@sycl@@YAX_N@Z ?ext_oneapi_get_state@queue@_V1@sycl@@QEBA?AW4queue_state@experimental@oneapi@ext@23@XZ ?ext_oneapi_graph@handler@_V1@sycl@@QEAAXV?$command_graph@$00@experimental@oneapi@ext@23@@Z ?ext_oneapi_graph@queue@_V1@sycl@@QEAA?AVevent@23@V?$command_graph@$00@experimental@oneapi@ext@23@AEBUcode_location@detail@23@@Z diff --git a/sycl/unittests/Extensions/CommandGraph.cpp b/sycl/unittests/Extensions/CommandGraph.cpp index 10cac8b18dcbd..c4ca2c93ab3fb 100644 --- a/sycl/unittests/Extensions/CommandGraph.cpp +++ b/sycl/unittests/Extensions/CommandGraph.cpp @@ -20,6 +20,24 @@ using namespace sycl; using namespace sycl::ext::oneapi; +// Spec constant for testing. +constexpr specialization_id SpecConst1{7}; + +namespace sycl { +inline namespace _V1 { +namespace detail { + +// Necessary for get_specialization_constant() to work in unit tests. +template <> const char *get_spec_constant_symbolic_ID() { + return "SC1"; +} +} // namespace detail +} // namespace _V1 +} // namespace sycl + +// anonymous namespace used to avoid code redundancy by defining functions +// used by multiple times by unitests. +// Defining anonymous namespace prevents from function naming conflits namespace { /// Define the three possible path to add node to a SYCL Graph. /// Shortcut is a sub-type of Record&Replay using Queue shortcut @@ -972,4 +990,70 @@ TEST_F(CommandGraphTest, Memcpy2DExceptionCheck) { sycl::free(USMMemSrc, Queue); sycl::free(USMMemDst, Queue); } ->>>>>>> b99238b99781... [SYCL][Graph] Implement exceptions for incompatible extensions (#276) + +// Tests that using specialization constants in a graph will throw. +TEST_F(CommandGraphTest, SpecializationConstant) { + + ASSERT_THROW( + { + try { + Graph.add([&](handler &CGH) { + CGH.set_specialization_constant(8); + }); + } catch (const sycl::exception &e) { + ASSERT_EQ(e.code(), make_error_code(sycl::errc::invalid)); + throw; + } + }, + sycl::exception); + ASSERT_THROW( + { + try { + Graph.add([&](handler &CGH) { + int Value = CGH.get_specialization_constant(); + (void)Value; + }); + } catch (const sycl::exception &e) { + ASSERT_EQ(e.code(), make_error_code(sycl::errc::invalid)); + throw; + } + }, + sycl::exception); +} + +// Tests that using kernel bundles in a graph will throw. +TEST_F(CommandGraphTest, KernelBundle) { + sycl::kernel_bundle KernelBundle = + sycl::get_kernel_bundle( + Queue.get_context(), {Dev}); + + ASSERT_THROW( + { + try { + Graph.add([&](handler &CGH) { CGH.use_kernel_bundle(KernelBundle); }); + } catch (const sycl::exception &e) { + ASSERT_EQ(e.code(), make_error_code(sycl::errc::invalid)); + throw; + } + }, + sycl::exception); +} + +// Tests that using reductions in a graph will throw. +TEST_F(CommandGraphTest, Reductions) { + int ReduVar = 0; + ASSERT_THROW( + { + try { + Graph.add([&](handler &CGH) { + CGH.parallel_for( + range<1>{1}, reduction(&ReduVar, int{0}, sycl::plus()), + [=](item<1> idx, auto &Sum) {}); + }); + } catch (const sycl::exception &e) { + ASSERT_EQ(e.code(), make_error_code(sycl::errc::invalid)); + throw; + } + }, + sycl::exception); +} From f1ffc70006c6a3dcfef6cd9a32fd3c3dc7414a18 Mon Sep 17 00:00:00 2001 From: Maxime France-Pillois Date: Thu, 10 Aug 2023 17:26:20 +0100 Subject: [PATCH 04/10] [SYCL][Graph] Throw exception when using bindless images extension in a graph (#294) * [SYCL][Graph] Throw exception when using bindless images extension in a graph Throws an invalid exception when using bindless images extension in a graph. Adds Unitests to test the exception throwing. --- .../sycl/ext/oneapi/experimental/graph.hpp | 5 +- sycl/source/handler.cpp | 25 +- sycl/unittests/Extensions/CommandGraph.cpp | 543 +++++++++++++++++- 3 files changed, 569 insertions(+), 4 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/graph.hpp b/sycl/include/sycl/ext/oneapi/experimental/graph.hpp index 3541cb62f1969..40f7c4cbbe51d 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/graph.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/graph.hpp @@ -40,7 +40,8 @@ enum class UnsupportedGraphFeatures { sycl_ext_oneapi_kernel_properties, sycl_ext_oneapi_enqueue_barrier, sycl_ext_oneapi_memcpy2d, - sycl_ext_oneapi_device_global + sycl_ext_oneapi_device_global, + sycl_ext_oneapi_bindless_images }; constexpr const char * @@ -61,6 +62,8 @@ UnsupportedFeatureToString(UnsupportedGraphFeatures Feature) { return "sycl_ext_oneapi_memcpy2d"; case UGF::sycl_ext_oneapi_device_global: return "sycl_ext_oneapi_device_global"; + case UGF::sycl_ext_oneapi_bindless_images: + return "sycl_ext_oneapi_bindless_images"; default: return {}; } diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index 69036e2dfa190..db64a6c141d3e 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -891,6 +891,9 @@ void handler::ext_oneapi_memset2d_impl(void *Dest, size_t DestPitch, int Value, void handler::ext_oneapi_copy( void *Src, ext::oneapi::experimental::image_mem_handle Dest, const ext::oneapi::experimental::image_descriptor &Desc) { + throwIfGraphAssociated< + ext::oneapi::experimental::detail::UnsupportedGraphFeatures:: + sycl_ext_oneapi_bindless_images>(); MSrcPtr = Src; MDstPtr = Dest.raw_handle; @@ -924,7 +927,9 @@ void handler::ext_oneapi_copy( ext::oneapi::experimental::image_mem_handle Dest, sycl::range<3> DestOffset, const ext::oneapi::experimental::image_descriptor &DestImgDesc, sycl::range<3> CopyExtent) { - + throwIfGraphAssociated< + ext::oneapi::experimental::detail::UnsupportedGraphFeatures:: + sycl_ext_oneapi_bindless_images>(); MSrcPtr = Src; MDstPtr = Dest.raw_handle; @@ -957,6 +962,9 @@ void handler::ext_oneapi_copy( void handler::ext_oneapi_copy( ext::oneapi::experimental::image_mem_handle Src, void *Dest, const ext::oneapi::experimental::image_descriptor &Desc) { + throwIfGraphAssociated< + ext::oneapi::experimental::detail::UnsupportedGraphFeatures:: + sycl_ext_oneapi_bindless_images>(); MSrcPtr = Src.raw_handle; MDstPtr = Dest; @@ -990,6 +998,9 @@ void handler::ext_oneapi_copy( const ext::oneapi::experimental::image_descriptor &SrcImgDesc, void *Dest, sycl::range<3> DestOffset, sycl::range<3> DestExtent, sycl::range<3> CopyExtent) { + throwIfGraphAssociated< + ext::oneapi::experimental::detail::UnsupportedGraphFeatures:: + sycl_ext_oneapi_bindless_images>(); MSrcPtr = Src.raw_handle; MDstPtr = Dest; @@ -1022,6 +1033,9 @@ void handler::ext_oneapi_copy( void handler::ext_oneapi_copy( void *Src, void *Dest, const ext::oneapi::experimental::image_descriptor &Desc, size_t Pitch) { + throwIfGraphAssociated< + ext::oneapi::experimental::detail::UnsupportedGraphFeatures:: + sycl_ext_oneapi_bindless_images>(); MSrcPtr = Src; MDstPtr = Dest; @@ -1057,6 +1071,9 @@ void handler::ext_oneapi_copy( const ext::oneapi::experimental::image_descriptor &DeviceImgDesc, size_t DeviceRowPitch, sycl::range<3> HostExtent, sycl::range<3> CopyExtent) { + throwIfGraphAssociated< + ext::oneapi::experimental::detail::UnsupportedGraphFeatures:: + sycl_ext_oneapi_bindless_images>(); MSrcPtr = Src; MDstPtr = Dest; @@ -1090,6 +1107,9 @@ void handler::ext_oneapi_copy( void handler::ext_oneapi_wait_external_semaphore( sycl::ext::oneapi::experimental::interop_semaphore_handle SemaphoreHandle) { + throwIfGraphAssociated< + ext::oneapi::experimental::detail::UnsupportedGraphFeatures:: + sycl_ext_oneapi_bindless_images>(); MImpl->MInteropSemaphoreHandle = (sycl::detail::pi::PiInteropSemaphoreHandle)SemaphoreHandle.raw_handle; setType(detail::CG::SemaphoreWait); @@ -1097,6 +1117,9 @@ void handler::ext_oneapi_wait_external_semaphore( void handler::ext_oneapi_signal_external_semaphore( sycl::ext::oneapi::experimental::interop_semaphore_handle SemaphoreHandle) { + throwIfGraphAssociated< + ext::oneapi::experimental::detail::UnsupportedGraphFeatures:: + sycl_ext_oneapi_bindless_images>(); MImpl->MInteropSemaphoreHandle = (sycl::detail::pi::PiInteropSemaphoreHandle)SemaphoreHandle.raw_handle; setType(detail::CG::SemaphoreSignal); diff --git a/sycl/unittests/Extensions/CommandGraph.cpp b/sycl/unittests/Extensions/CommandGraph.cpp index c4ca2c93ab3fb..30055a60e3715 100644 --- a/sycl/unittests/Extensions/CommandGraph.cpp +++ b/sycl/unittests/Extensions/CommandGraph.cpp @@ -309,7 +309,245 @@ void addMemcpy2D(experimental::detail::modifiable_command_graph &G, queue &Q, ASSERT_EQ(ExceptionCode, sycl::errc::invalid); } -} // namespace +/// Tries to add nodes including images bindless copy instructions +/// to the graph G. It tests that an invalid exception has been thrown +/// Since sycl_ext_oneapi_bindless_images extension can not be used +/// along with SYCL Graph. +/// +/// @param G Modifiable graph to add commands to. +/// @param Q Queue to submit nodes to. +/// @param Img Image memory +/// @param HostData Host Pointer to the memory +/// @param ImgUSM USM Pointer to Image memory +/// @param Pitch image pitch +/// @param Desc Image descriptor +template +void addImagesCopies(experimental::detail::modifiable_command_graph &G, + queue &Q, sycl::ext::oneapi::experimental::image_mem Img, + std::vector HostData, void *ImgUSM, + size_t Pitch, + sycl::ext::oneapi::experimental::image_descriptor Desc) { + // simple copy Host to Device + std::error_code ExceptionCode = make_error_code(sycl::errc::success); + try { + if constexpr (PathKind == OperationPath::RecordReplay) { + Q.submit([&](handler &CGH) { + CGH.ext_oneapi_copy(HostData.data(), Img.get_handle(), Desc); + }); + } + if constexpr (PathKind == OperationPath::Shortcut) { + Q.ext_oneapi_copy(HostData.data(), Img.get_handle(), Desc); + } + if constexpr (PathKind == OperationPath::Explicit) { + G.add([&](handler &CGH) { + CGH.ext_oneapi_copy(HostData.data(), Img.get_handle(), Desc); + }); + } + } catch (exception &Exception) { + ExceptionCode = Exception.code(); + } + ASSERT_EQ(ExceptionCode, sycl::errc::invalid); + + // simple copy Device to Host + ExceptionCode = make_error_code(sycl::errc::success); + try { + if constexpr (PathKind == OperationPath::RecordReplay) { + Q.submit([&](handler &CGH) { + CGH.ext_oneapi_copy(Img.get_handle(), HostData.data(), Desc); + }); + } + if constexpr (PathKind == OperationPath::Shortcut) { + Q.ext_oneapi_copy(Img.get_handle(), HostData.data(), Desc); + } + if constexpr (PathKind == OperationPath::Explicit) { + G.add([&](handler &CGH) { + CGH.ext_oneapi_copy(Img.get_handle(), HostData.data(), Desc); + }); + } + } catch (exception &Exception) { + ExceptionCode = Exception.code(); + } + ASSERT_EQ(ExceptionCode, sycl::errc::invalid); + + // simple copy Host to Device USM + ExceptionCode = make_error_code(sycl::errc::success); + try { + if constexpr (PathKind == OperationPath::RecordReplay) { + Q.submit([&](handler &CGH) { + CGH.ext_oneapi_copy(HostData.data(), ImgUSM, Desc, Pitch); + }); + } + if constexpr (PathKind == OperationPath::Shortcut) { + Q.ext_oneapi_copy(HostData.data(), ImgUSM, Desc, Pitch); + } + if constexpr (PathKind == OperationPath::Explicit) { + G.add([&](handler &CGH) { + CGH.ext_oneapi_copy(HostData.data(), ImgUSM, Desc, Pitch); + }); + } + } catch (exception &Exception) { + ExceptionCode = Exception.code(); + } + ASSERT_EQ(ExceptionCode, sycl::errc::invalid); + + // subregion copy Host to Device + ExceptionCode = make_error_code(sycl::errc::success); + try { + if constexpr (PathKind == OperationPath::RecordReplay) { + Q.submit([&](handler &CGH) { + CGH.ext_oneapi_copy(HostData.data(), {0, 0, 0}, {0, 0, 0}, + Img.get_handle(), {0, 0, 0}, Desc, {0, 0, 0}); + }); + } + if constexpr (PathKind == OperationPath::Shortcut) { + Q.ext_oneapi_copy(HostData.data(), {0, 0, 0}, {0, 0, 0}, Img.get_handle(), + {0, 0, 0}, Desc, {0, 0, 0}); + } + if constexpr (PathKind == OperationPath::Explicit) { + G.add([&](handler &CGH) { + CGH.ext_oneapi_copy(HostData.data(), {0, 0, 0}, {0, 0, 0}, + Img.get_handle(), {0, 0, 0}, Desc, {0, 0, 0}); + }); + } + } catch (exception &Exception) { + ExceptionCode = Exception.code(); + } + ASSERT_EQ(ExceptionCode, sycl::errc::invalid); + + // subregion copy Device to Host + ExceptionCode = make_error_code(sycl::errc::success); + try { + if constexpr (PathKind == OperationPath::RecordReplay) { + Q.submit([&](handler &CGH) { + CGH.ext_oneapi_copy(Img.get_handle(), {0, 0, 0}, Desc, HostData.data(), + {0, 0, 0}, {0, 0, 0}, {0, 0, 0}); + }); + } + if constexpr (PathKind == OperationPath::Shortcut) { + Q.ext_oneapi_copy(Img.get_handle(), {0, 0, 0}, Desc, HostData.data(), + {0, 0, 0}, {0, 0, 0}, {0, 0, 0}); + } + if constexpr (PathKind == OperationPath::Explicit) { + G.add([&](handler &CGH) { + CGH.ext_oneapi_copy(Img.get_handle(), {0, 0, 0}, Desc, HostData.data(), + {0, 0, 0}, {0, 0, 0}, {0, 0, 0}); + }); + } + } catch (exception &Exception) { + ExceptionCode = Exception.code(); + } + ASSERT_EQ(ExceptionCode, sycl::errc::invalid); + + // subregion copy Host to Device USM + ExceptionCode = make_error_code(sycl::errc::success); + try { + if constexpr (PathKind == OperationPath::RecordReplay) { + Q.submit([&](handler &CGH) { + CGH.ext_oneapi_copy(HostData.data(), {0, 0, 0}, ImgUSM, {0, 0, 0}, Desc, + Pitch, {0, 0, 0}, {0, 0, 0}); + }); + } + if constexpr (PathKind == OperationPath::Shortcut) { + Q.ext_oneapi_copy(HostData.data(), {0, 0, 0}, ImgUSM, {0, 0, 0}, Desc, + Pitch, {0, 0, 0}, {0, 0, 0}); + } + if constexpr (PathKind == OperationPath::Explicit) { + G.add([&](handler &CGH) { + CGH.ext_oneapi_copy(HostData.data(), {0, 0, 0}, ImgUSM, {0, 0, 0}, Desc, + Pitch, {0, 0, 0}, {0, 0, 0}); + }); + } + } catch (exception &Exception) { + ExceptionCode = Exception.code(); + } + ASSERT_EQ(ExceptionCode, sycl::errc::invalid); +} + +bool depthSearchSuccessorCheck( + std::shared_ptr Node) { + if (Node->MSuccessors.size() > 1) + return false; + + for (const auto &Succ : Node->MSuccessors) { + return Succ->depthSearchCount(); + } + return true; +} + +/// Submits four kernels with diamond dependency to the queue Q +/// @param Q Queue to submit nodes to. +void runKernels(queue Q) { + auto NodeA = Q.submit( + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); + auto NodeB = Q.submit([&](sycl::handler &cgh) { + cgh.depends_on(NodeA); + cgh.single_task>([]() {}); + }); + auto NodeC = Q.submit([&](sycl::handler &cgh) { + cgh.depends_on(NodeA); + cgh.single_task>([]() {}); + }); + auto NodeD = Q.submit([&](sycl::handler &cgh) { + cgh.depends_on({NodeB, NodeC}); + cgh.single_task>([]() {}); + }); +} + +/// Submits four kernels without any additional dependencies the queue Q +/// @param Q Queue to submit nodes to. +void runKernelsInOrder(queue Q) { + auto NodeA = Q.submit( + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); + auto NodeB = Q.submit( + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); + auto NodeC = Q.submit( + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); + auto NodeD = Q.submit( + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); +} + +/// Adds four kernels with diamond dependency to the Graph G +/// @param G Modifiable graph to add commands to. +void addKernels( + experimental::command_graph G) { + auto NodeA = G.add( + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); + auto NodeB = + G.add([&](sycl::handler &cgh) { cgh.single_task>([]() {}); }, + {experimental::property::node::depends_on(NodeA)}); + auto NodeC = + G.add([&](sycl::handler &cgh) { cgh.single_task>([]() {}); }, + {experimental::property::node::depends_on(NodeA)}); + auto NodeD = + G.add([&](sycl::handler &cgh) { cgh.single_task>([]() {}); }, + {experimental::property::node::depends_on(NodeB, NodeC)}); +} + +bool checkExecGraphSchedule( + std::shared_ptr + GraphA, + std::shared_ptr + GraphB) { + auto ScheduleA = GraphA->getSchedule(); + auto ScheduleB = GraphB->getSchedule(); + if (ScheduleA.size() != ScheduleB.size()) + return false; + + std::vector< + std::shared_ptr> + VScheduleA{std::begin(ScheduleA), std::end(ScheduleA)}; + std::vector< + std::shared_ptr> + VScheduleB{std::begin(ScheduleB), std::end(ScheduleB)}; + + for (size_t i = 0; i < VScheduleA.size(); i++) { + if (!VScheduleA[i]->isSimilar(VScheduleB[i])) + return false; + } + return true; +} + +} // anonymous namespace class CommandGraphTest : public ::testing::Test { public: @@ -1046,7 +1284,7 @@ TEST_F(CommandGraphTest, Reductions) { { try { Graph.add([&](handler &CGH) { - CGH.parallel_for( + CGH.parallel_for( range<1>{1}, reduction(&ReduVar, int{0}, sycl::plus()), [=](item<1> idx, auto &Sum) {}); }); @@ -1057,3 +1295,304 @@ TEST_F(CommandGraphTest, Reductions) { }, sycl::exception); } + +TEST_F(CommandGraphTest, BindlessExceptionCheck) { + auto Ctxt = Queue.get_context(); + + // declare image data + size_t Height = 13; + size_t Width = 7; + size_t Depth = 11; + size_t N = Height * Width * Depth; + std::vector DataIn(N); + + // Extension: image descriptor - can use the same for both images + sycl::ext::oneapi::experimental::image_descriptor Desc( + {Width, Height, Depth}, sycl::image_channel_order::rgba, + sycl::image_channel_type::fp32); + + // Extension: allocate memory on device and create the handle + // Input images memory + sycl::ext::oneapi::experimental::image_mem ImgMem(Desc, Dev, Ctxt); + // Extension: returns the device pointer to USM allocated pitched memory + size_t Pitch = 0; + auto ImgMemUSM = sycl::ext::oneapi::experimental::pitched_alloc_device( + &Pitch, Desc, Queue); + + Graph.begin_recording(Queue); + + addImagesCopies(Graph, Queue, ImgMem, DataIn, + ImgMemUSM, Pitch, Desc); + + addImagesCopies(Graph, Queue, ImgMem, DataIn, + ImgMemUSM, Pitch, Desc); + + Graph.end_recording(); + + addImagesCopies(Graph, Queue, ImgMem, DataIn, + ImgMemUSM, Pitch, Desc); + + sycl::free(ImgMemUSM, Ctxt); +} + +class MultiThreadGraphTest : public CommandGraphTest { +public: + MultiThreadGraphTest() + : CommandGraphTest(), NumThreads(std::thread::hardware_concurrency()), + SyncPoint(NumThreads) { + Threads.reserve(NumThreads); + } + +protected: + const unsigned NumThreads; + Barrier SyncPoint; + std::vector Threads; +}; + +TEST_F(MultiThreadGraphTest, BeginEndRecording) { + auto RecordGraph = [&]() { + queue MyQueue{Queue.get_context(), Queue.get_device()}; + + SyncPoint.wait(); + + Graph.begin_recording(MyQueue); + runKernels(MyQueue); + Graph.end_recording(MyQueue); + }; + + for (unsigned i = 0; i < NumThreads; ++i) { + Threads.emplace_back(RecordGraph); + } + + for (unsigned i = 0; i < NumThreads; ++i) { + Threads[i].join(); + } + + // Reference computation + queue QueueRef; + experimental::command_graph GraphRef{ + Queue.get_context(), Queue.get_device()}; + + for (unsigned i = 0; i < NumThreads; ++i) { + queue MyQueue; + GraphRef.begin_recording(MyQueue); + runKernels(MyQueue); + GraphRef.end_recording(MyQueue); + } + + auto GraphImpl = sycl::detail::getSyclObjImpl(Graph); + auto GraphRefImpl = sycl::detail::getSyclObjImpl(GraphRef); + ASSERT_EQ(GraphImpl->hasSimilarStructure(GraphRefImpl), true); +} + +TEST_F(MultiThreadGraphTest, ExplicitAddNodes) { + auto RecordGraph = [&]() { + queue MyQueue{Queue.get_context(), Queue.get_device()}; + + SyncPoint.wait(); + addKernels(Graph); + }; + + for (unsigned i = 0; i < NumThreads; ++i) { + Threads.emplace_back(RecordGraph); + } + + for (unsigned i = 0; i < NumThreads; ++i) { + Threads[i].join(); + } + + // Reference computation + queue QueueRef; + experimental::command_graph GraphRef{ + Queue.get_context(), Queue.get_device()}; + + for (unsigned i = 0; i < NumThreads; ++i) { + addKernels(GraphRef); + } + + auto GraphImpl = sycl::detail::getSyclObjImpl(Graph); + auto GraphRefImpl = sycl::detail::getSyclObjImpl(GraphRef); + ASSERT_EQ(GraphImpl->hasSimilarStructure(GraphRefImpl), true); +} + +TEST_F(MultiThreadGraphTest, RecordAddNodes) { + Graph.begin_recording(Queue); + auto RecordGraph = [&]() { + SyncPoint.wait(); + runKernels(Queue); + }; + + for (unsigned i = 0; i < NumThreads; ++i) { + Threads.emplace_back(RecordGraph); + } + + for (unsigned i = 0; i < NumThreads; ++i) { + Threads[i].join(); + } + + // We stop recording the Queue when all threads have finished their processing + Graph.end_recording(Queue); + + // Reference computation + queue QueueRef; + experimental::command_graph GraphRef{ + Queue.get_context(), Queue.get_device()}; + + GraphRef.begin_recording(QueueRef); + for (unsigned i = 0; i < NumThreads; ++i) { + runKernels(QueueRef); + } + GraphRef.end_recording(QueueRef); + + auto GraphImpl = sycl::detail::getSyclObjImpl(Graph); + auto GraphRefImpl = sycl::detail::getSyclObjImpl(GraphRef); + ASSERT_EQ(GraphImpl->hasSimilarStructure(GraphRefImpl), true); +} + +TEST_F(MultiThreadGraphTest, RecordAddNodesInOrderQueue) { + sycl::property_list Properties{sycl::property::queue::in_order()}; + queue InOrderQueue{Dev, Properties}; + + experimental::command_graph + InOrderGraph{InOrderQueue.get_context(), InOrderQueue.get_device()}; + + InOrderGraph.begin_recording(InOrderQueue); + auto RecordGraph = [&]() { + SyncPoint.wait(); + runKernelsInOrder(InOrderQueue); + }; + + for (unsigned i = 0; i < NumThreads; ++i) { + Threads.emplace_back(RecordGraph); + } + + for (unsigned i = 0; i < NumThreads; ++i) { + Threads[i].join(); + } + + // We stop recording the Queue when all threads have finished their processing + InOrderGraph.end_recording(InOrderQueue); + + // Reference computation + queue InOrderQueueRef{Dev, Properties}; + experimental::command_graph + InOrderGraphRef{InOrderQueueRef.get_context(), + InOrderQueueRef.get_device()}; + + InOrderGraphRef.begin_recording(InOrderQueueRef); + for (unsigned i = 0; i < NumThreads; ++i) { + runKernelsInOrder(InOrderQueueRef); + } + InOrderGraphRef.end_recording(InOrderQueueRef); + + auto GraphImpl = sycl::detail::getSyclObjImpl(InOrderGraph); + auto GraphRefImpl = sycl::detail::getSyclObjImpl(InOrderGraphRef); + ASSERT_EQ(GraphImpl->getNumberOfNodes(), GraphRefImpl->getNumberOfNodes()); + + // In-order graph must have only a single root + ASSERT_EQ(GraphImpl->MRoots.size(), 1lu); + + // Check structure graph + for (auto Node : GraphImpl->MRoots) { + ASSERT_EQ(depthSearchSuccessorCheck(Node), true); + } +} + +TEST_F(MultiThreadGraphTest, Finalize) { + addKernels(Graph); + + std::map> + GraphsExecMap; + auto FinalizeGraph = [&](int ThreadNum) { + SyncPoint.wait(); + auto GraphExec = Graph.finalize(); + + GraphsExecMap.insert( + std::map>:: + value_type(ThreadNum, GraphExec)); + Queue.submit([&](sycl::handler &CGH) { CGH.ext_oneapi_graph(GraphExec); }); + }; + + for (unsigned i = 0; i < NumThreads; ++i) { + Threads.emplace_back(FinalizeGraph, i); + } + + for (unsigned i = 0; i < NumThreads; ++i) { + Threads[i].join(); + } + + // Reference computation + queue QueueRef; + experimental::command_graph GraphRef{ + Queue.get_context(), Queue.get_device()}; + + addKernels(GraphRef); + + for (unsigned i = 0; i < NumThreads; ++i) { + auto GraphExecRef = GraphRef.finalize(); + QueueRef.submit( + [&](sycl::handler &CGH) { CGH.ext_oneapi_graph(GraphExecRef); }); + auto GraphExecImpl = + sycl::detail::getSyclObjImpl(GraphsExecMap.find(i)->second); + auto GraphExecRefImpl = sycl::detail::getSyclObjImpl(GraphExecRef); + ASSERT_EQ(checkExecGraphSchedule(GraphExecImpl, GraphExecRefImpl), true); + } +} + +TEST_F(CommandGraphTest, InvalidBuffer) { + // Check that using a buffer with write_back enabled in a graph will throw. + int Data; + // Create a buffer which does not have write-back disabled. + buffer Buffer{&Data, range<1>{1}}; + + // Use this buffer in the graph, this should throw. + ASSERT_THROW( + { + try { + Graph.add([&](handler &CGH) { + auto Acc = Buffer.get_access(CGH); + }); + } catch (const sycl::exception &e) { + ASSERT_EQ(e.code(), make_error_code(sycl::errc::invalid)); + throw; + } + }, + sycl::exception); +} + +TEST_F(CommandGraphTest, InvalidHostAccessor) { + // Check that creating a host_accessor on a buffer which is in use by a graph + // will throw. + + // Create a buffer which does not have write-back disabled. + buffer Buffer{range<1>{1}}; + + { + // Create a graph in local scope so we can destroy it + ext::oneapi::experimental::command_graph Graph{ + Queue.get_context(), + Queue.get_device(), + {experimental::property::graph::assume_buffer_outlives_graph{}}}; + + // Add the buffer to the graph. + Graph.add([&](handler &CGH) { + auto Acc = Buffer.get_access(CGH); + }); + + // Attempt to create a host_accessor, which should throw. + ASSERT_THROW( + { + try { + host_accessor HostAcc{Buffer}; + } catch (const sycl::exception &e) { + ASSERT_EQ(e.code(), make_error_code(sycl::errc::invalid)); + throw; + } + }, + sycl::exception); + } + // Graph is now out of scope so we should be able to create a host_accessor + ASSERT_NO_THROW({ host_accessor HostAcc{Buffer}; }); +} From 5c37e77892d54cd16b38896b060815cca681ef36 Mon Sep 17 00:00:00 2001 From: Maxime France-Pillois Date: Thu, 10 Aug 2023 17:46:43 +0100 Subject: [PATCH 05/10] [SYCL][Graph] Throw exception when calling profiling (#295) Throws an invalid exception when calling `event::get_profiling_info()` on an event returned by a graph submission. Adds unitests checking that exceptions are throws if `event::get_profiling_info()` is called on an event returned from a queue in recording mode or an event returned by a graph submission. --- sycl/source/detail/event_impl.cpp | 5 ++ sycl/source/detail/event_impl.hpp | 10 ++++ sycl/source/detail/graph_impl.cpp | 1 + sycl/unittests/Extensions/CommandGraph.cpp | 70 ++++++++++++++++++++++ 4 files changed, 86 insertions(+) diff --git a/sycl/source/detail/event_impl.cpp b/sycl/source/detail/event_impl.cpp index 85893d0c1a14d..21a1e720dd070 100644 --- a/sycl/source/detail/event_impl.cpp +++ b/sycl/source/detail/event_impl.cpp @@ -274,6 +274,11 @@ void event_impl::checkProfilingPreconditions() const { "Profiling information is unavailable as the queue associated with " "the event does not have the 'enable_profiling' property."); } + if (EventFromSubmitedExecCommandBuffer) { + throw sycl::exception(make_error_code(sycl::errc::invalid), + "Profiling information is unavailable for events " + "returned by a graph submission."); + } } template <> diff --git a/sycl/source/detail/event_impl.hpp b/sycl/source/detail/event_impl.hpp index 4259ab803b3be..2bcad45bd9d33 100644 --- a/sycl/source/detail/event_impl.hpp +++ b/sycl/source/detail/event_impl.hpp @@ -278,6 +278,14 @@ class event_impl { return MGraph.lock(); } + void setEventFromSubmitedExecCommandBuffer(bool value) { + EventFromSubmitedExecCommandBuffer = value; + } + + bool isEventFromSubmitedExecCommandBuffer() const { + return EventFromSubmitedExecCommandBuffer; + } + protected: // When instrumentation is enabled emits trace event for event wait begin and // returns the telemetry event generated for the wait @@ -327,6 +335,8 @@ class event_impl { /// Store the command graph associated with this event, if any. /// This event is also be stored in the graph so a weak_ptr is used. std::weak_ptr MGraph; + /// Indicates that the event results from a command graph submission + bool EventFromSubmitedExecCommandBuffer = false; // If this event represents a submission to a // sycl::detail::pi::PiExtCommandBuffer the sync point for that submission is diff --git a/sycl/source/detail/graph_impl.cpp b/sycl/source/detail/graph_impl.cpp index 143172c3ea80a..47bb5921d25eb 100644 --- a/sycl/source/detail/graph_impl.cpp +++ b/sycl/source/detail/graph_impl.cpp @@ -372,6 +372,7 @@ exec_graph_impl::enqueue(const std::shared_ptr &Queue, auto NewEvent = std::make_shared(Queue); NewEvent->setContextImpl(Queue->getContextImplPtr()); NewEvent->setStateIncomplete(); + NewEvent->setEventFromSubmitedExecCommandBuffer(true); return NewEvent; }); diff --git a/sycl/unittests/Extensions/CommandGraph.cpp b/sycl/unittests/Extensions/CommandGraph.cpp index 30055a60e3715..30a0b63f676d7 100644 --- a/sycl/unittests/Extensions/CommandGraph.cpp +++ b/sycl/unittests/Extensions/CommandGraph.cpp @@ -1335,6 +1335,76 @@ TEST_F(CommandGraphTest, BindlessExceptionCheck) { sycl::free(ImgMemUSM, Ctxt); } +TEST_F(CommandGraphTest, GetProfilingInfoExceptionCheck) { + sycl::context Ctx{Dev}; + sycl::queue QueueProfile{ + Ctx, Dev, sycl::property_list{sycl::property::queue::enable_profiling{}}}; + experimental::command_graph + GraphProfile{QueueProfile.get_context(), Dev}; + + GraphProfile.begin_recording(QueueProfile); + auto Event = QueueProfile.submit( + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); + + // Checks that exception is thrown when get_profile_info is called on "event" + // returned by a queue in recording mode. + std::error_code ExceptionCode = make_error_code(sycl::errc::success); + try { + Event.get_profiling_info(); + } catch (exception &Exception) { + ExceptionCode = Exception.code(); + } + ASSERT_EQ(ExceptionCode, sycl::errc::invalid); + + ExceptionCode = make_error_code(sycl::errc::success); + try { + Event.get_profiling_info(); + } catch (exception &Exception) { + ExceptionCode = Exception.code(); + } + ASSERT_EQ(ExceptionCode, sycl::errc::invalid); + + ExceptionCode = make_error_code(sycl::errc::success); + try { + Event.get_profiling_info(); + } catch (exception &Exception) { + ExceptionCode = Exception.code(); + } + ASSERT_EQ(ExceptionCode, sycl::errc::invalid); + + GraphProfile.end_recording(); + + auto GraphExec = GraphProfile.finalize(); + auto EventSub = QueueProfile.submit( + [&](sycl::handler &CGH) { CGH.ext_oneapi_graph(GraphExec); }); + + // Checks that exception is thrown when get_profile_info is called on "event" + // returned by a graph submission. + ExceptionCode = make_error_code(sycl::errc::success); + try { + EventSub.get_profiling_info(); + } catch (exception &Exception) { + ExceptionCode = Exception.code(); + } + ASSERT_EQ(ExceptionCode, sycl::errc::invalid); + + ExceptionCode = make_error_code(sycl::errc::success); + try { + EventSub.get_profiling_info(); + } catch (exception &Exception) { + ExceptionCode = Exception.code(); + } + ASSERT_EQ(ExceptionCode, sycl::errc::invalid); + + ExceptionCode = make_error_code(sycl::errc::success); + try { + EventSub.get_profiling_info(); + } catch (exception &Exception) { + ExceptionCode = Exception.code(); + } + ASSERT_EQ(ExceptionCode, sycl::errc::invalid); +} + class MultiThreadGraphTest : public CommandGraphTest { public: MultiThreadGraphTest() From 17887faf1e6564b1649b8460aa86f71982268f0f Mon Sep 17 00:00:00 2001 From: Ewan Crawford Date: Fri, 11 Aug 2023 13:17:32 +0100 Subject: [PATCH 06/10] Fixup cherry-pick issues --- sycl/test/abi/sycl_symbols_linux.dump | 1 + sycl/test/abi/sycl_symbols_windows.dump | 2 +- sycl/unittests/Extensions/CommandGraph.cpp | 396 ++------------------- sycl/unittests/helpers/PiMockPlugin.hpp | 2 +- 4 files changed, 28 insertions(+), 373 deletions(-) diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 1967793fea90a..4f18f49fbb098 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -4674,6 +4674,7 @@ _ZNK4sycl3_V17context8get_infoINS0_4info7context8platformEEENS0_6detail20is_cont _ZNK4sycl3_V17context9getNativeEv _ZNK4sycl3_V17handler15getCommandGraphEv _ZNK4sycl3_V17handler17getContextImplPtrEv +_ZNK4sycl3_V17handler22throwIfGraphAssociatedILNS0_3ext6oneapi12experimental6detail24UnsupportedGraphFeaturesE7EEEvv _ZNK4sycl3_V17handler27isStateExplicitKernelBundleEv _ZNK4sycl3_V17handler30getOrInsertHandlerKernelBundleEb _ZNK4sycl3_V17sampler12get_propertyINS0_3ext6oneapi4cuda8property7context19use_primary_contextEEET_v diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index d51c06d332cb0..545bdb4361454 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -979,7 +979,6 @@ ?die@pi@detail@_V1@sycl@@YAXPEBD@Z ?discard_or_return@queue@_V1@sycl@@AEAA?AVevent@23@AEBV423@@Z ?empty@kernel_bundle_plain@detail@_V1@sycl@@QEBA_NXZ -?enable_ext_oneapi_default_context@detail@_V1@sycl@@YAX_N@Z ?end@HostProfilingInfo@detail@_V1@sycl@@QEAAXXZ ?end@exception_list@_V1@sycl@@QEBA?AV?$_Vector_const_iterator@V?$_Vector_val@U?$_Simple_types@Vexception_ptr@std@@@std@@@std@@@std@@XZ ?end@kernel_bundle_plain@detail@_V1@sycl@@IEBAPEBVdevice_image_plain@234@XZ @@ -1026,6 +1025,7 @@ ?ext_oneapi_enable_peer_access@device@_V1@sycl@@QEAAXAEBV123@@Z ?ext_oneapi_fill2d_impl@handler@_V1@sycl@@AEAAXPEAX_KPEBX111@Z ?ext_oneapi_get_default_context@platform@_V1@sycl@@QEBA?AVcontext@23@XZ +?enable_ext_oneapi_default_context@detail@_V1@sycl@@YAX_N@Z ?ext_oneapi_get_state@queue@_V1@sycl@@QEBA?AW4queue_state@experimental@oneapi@ext@23@XZ ?ext_oneapi_graph@handler@_V1@sycl@@QEAAXV?$command_graph@$00@experimental@oneapi@ext@23@@Z ?ext_oneapi_graph@queue@_V1@sycl@@QEAA?AVevent@23@V?$command_graph@$00@experimental@oneapi@ext@23@AEBUcode_location@detail@23@@Z diff --git a/sycl/unittests/Extensions/CommandGraph.cpp b/sycl/unittests/Extensions/CommandGraph.cpp index 30a0b63f676d7..2a8eb0242b66a 100644 --- a/sycl/unittests/Extensions/CommandGraph.cpp +++ b/sycl/unittests/Extensions/CommandGraph.cpp @@ -14,6 +14,7 @@ #include #include #include +#include #include @@ -462,91 +463,6 @@ void addImagesCopies(experimental::detail::modifiable_command_graph &G, } ASSERT_EQ(ExceptionCode, sycl::errc::invalid); } - -bool depthSearchSuccessorCheck( - std::shared_ptr Node) { - if (Node->MSuccessors.size() > 1) - return false; - - for (const auto &Succ : Node->MSuccessors) { - return Succ->depthSearchCount(); - } - return true; -} - -/// Submits four kernels with diamond dependency to the queue Q -/// @param Q Queue to submit nodes to. -void runKernels(queue Q) { - auto NodeA = Q.submit( - [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); - auto NodeB = Q.submit([&](sycl::handler &cgh) { - cgh.depends_on(NodeA); - cgh.single_task>([]() {}); - }); - auto NodeC = Q.submit([&](sycl::handler &cgh) { - cgh.depends_on(NodeA); - cgh.single_task>([]() {}); - }); - auto NodeD = Q.submit([&](sycl::handler &cgh) { - cgh.depends_on({NodeB, NodeC}); - cgh.single_task>([]() {}); - }); -} - -/// Submits four kernels without any additional dependencies the queue Q -/// @param Q Queue to submit nodes to. -void runKernelsInOrder(queue Q) { - auto NodeA = Q.submit( - [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); - auto NodeB = Q.submit( - [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); - auto NodeC = Q.submit( - [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); - auto NodeD = Q.submit( - [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); -} - -/// Adds four kernels with diamond dependency to the Graph G -/// @param G Modifiable graph to add commands to. -void addKernels( - experimental::command_graph G) { - auto NodeA = G.add( - [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); - auto NodeB = - G.add([&](sycl::handler &cgh) { cgh.single_task>([]() {}); }, - {experimental::property::node::depends_on(NodeA)}); - auto NodeC = - G.add([&](sycl::handler &cgh) { cgh.single_task>([]() {}); }, - {experimental::property::node::depends_on(NodeA)}); - auto NodeD = - G.add([&](sycl::handler &cgh) { cgh.single_task>([]() {}); }, - {experimental::property::node::depends_on(NodeB, NodeC)}); -} - -bool checkExecGraphSchedule( - std::shared_ptr - GraphA, - std::shared_ptr - GraphB) { - auto ScheduleA = GraphA->getSchedule(); - auto ScheduleB = GraphB->getSchedule(); - if (ScheduleA.size() != ScheduleB.size()) - return false; - - std::vector< - std::shared_ptr> - VScheduleA{std::begin(ScheduleA), std::end(ScheduleA)}; - std::vector< - std::shared_ptr> - VScheduleB{std::begin(ScheduleB), std::end(ScheduleB)}; - - for (size_t i = 0; i < VScheduleA.size(); i++) { - if (!VScheduleA[i]->isSimilar(VScheduleB[i])) - return false; - } - return true; -} - } // anonymous namespace class CommandGraphTest : public ::testing::Test { @@ -586,7 +502,7 @@ TEST_F(CommandGraphTest, AddNode) { ASSERT_TRUE(GraphImpl->MRoots.empty()); auto Node1 = Graph.add( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); ASSERT_NE(sycl::detail::getSyclObjImpl(Node1), nullptr); ASSERT_FALSE(sycl::detail::getSyclObjImpl(Node1)->isEmpty()); ASSERT_EQ(GraphImpl->MRoots.size(), 1lu); @@ -641,17 +557,17 @@ TEST_F(CommandGraphTest, Finalize) { sycl::buffer Buf(1); auto Node1 = Graph.add([&](sycl::handler &cgh) { sycl::accessor A(Buf, cgh, sycl::write_only, sycl::no_init); - cgh.single_task([=]() { A[0] = 1; }); + cgh.single_task>([]() {}); }); // Add independent node auto Node2 = Graph.add( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); // Add a node that depends on Node1 due to the accessor auto Node3 = Graph.add([&](sycl::handler &cgh) { sycl::accessor A(Buf, cgh, sycl::write_only, sycl::no_init); - cgh.single_task([=]() { A[0] = 3; }); + cgh.single_task>([]() {}); }); // Guarantee order of independent nodes 1 and 2 @@ -677,7 +593,7 @@ TEST_F(CommandGraphTest, MakeEdge) { // Add two independent nodes auto Node1 = Graph.add( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); auto Node2 = Graph.add([&](sycl::handler &cgh) {}); ASSERT_EQ(GraphImpl->MRoots.size(), 2ul); ASSERT_TRUE(sycl::detail::getSyclObjImpl(Node1)->MSuccessors.empty()); @@ -771,7 +687,7 @@ TEST_F(CommandGraphTest, BeginEndRecording) { TEST_F(CommandGraphTest, GetCGCopy) { auto Node1 = Graph.add([&](sycl::handler &cgh) {}); auto Node2 = Graph.add( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }, + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }, {experimental::property::node::depends_on(Node1)}); // Get copy of CG of Node2 and check equality @@ -793,21 +709,21 @@ TEST_F(CommandGraphTest, GetCGCopy) { TEST_F(CommandGraphTest, SubGraph) { // Add sub-graph with two nodes auto Node1Graph = Graph.add( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); auto Node2Graph = Graph.add( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }, + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }, {experimental::property::node::depends_on(Node1Graph)}); auto GraphExec = Graph.finalize(); // Add node to main graph followed by sub-graph and another node experimental::command_graph MainGraph(Queue.get_context(), Dev); auto Node1MainGraph = MainGraph.add( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); auto Node2MainGraph = MainGraph.add([&](handler &CGH) { CGH.ext_oneapi_graph(GraphExec); }, {experimental::property::node::depends_on(Node1MainGraph)}); auto Node3MainGraph = MainGraph.add( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }, + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }, {experimental::property::node::depends_on(Node2MainGraph)}); // Assert order of the added sub-graph @@ -845,10 +761,10 @@ TEST_F(CommandGraphTest, RecordSubGraph) { // Record sub-graph with two nodes Graph.begin_recording(Queue); auto Node1Graph = Queue.submit( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); auto Node2Graph = Queue.submit([&](sycl::handler &cgh) { cgh.depends_on(Node1Graph); - cgh.single_task([]() {}); + cgh.single_task>([]() {}); }); Graph.end_recording(Queue); auto GraphExec = Graph.finalize(); @@ -857,14 +773,14 @@ TEST_F(CommandGraphTest, RecordSubGraph) { experimental::command_graph MainGraph(Queue.get_context(), Dev); MainGraph.begin_recording(Queue); auto Node1MainGraph = Queue.submit( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); auto Node2MainGraph = Queue.submit([&](handler &cgh) { cgh.depends_on(Node1MainGraph); cgh.ext_oneapi_graph(GraphExec); }); auto Node3MainGraph = Queue.submit([&](sycl::handler &cgh) { cgh.depends_on(Node2MainGraph); - cgh.single_task([]() {}); + cgh.single_task>([]() {}); }); MainGraph.end_recording(Queue); @@ -914,7 +830,7 @@ TEST_F(CommandGraphTest, InOrderQueue) { // Record in-order queue with three nodes InOrderGraph.begin_recording(InOrderQueue); auto Node1Graph = InOrderQueue.submit( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); auto PtrNode1 = sycl::detail::getSyclObjImpl(InOrderGraph) @@ -923,7 +839,7 @@ TEST_F(CommandGraphTest, InOrderQueue) { ASSERT_TRUE(PtrNode1->MPredecessors.empty()); auto Node2Graph = InOrderQueue.submit( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); auto PtrNode2 = sycl::detail::getSyclObjImpl(InOrderGraph) @@ -936,7 +852,7 @@ TEST_F(CommandGraphTest, InOrderQueue) { ASSERT_EQ(PtrNode2->MPredecessors.front().lock(), PtrNode1); auto Node3Graph = InOrderQueue.submit( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); auto PtrNode3 = sycl::detail::getSyclObjImpl(InOrderGraph) @@ -974,7 +890,7 @@ TEST_F(CommandGraphTest, InOrderQueueWithEmpty) { // node InOrderGraph.begin_recording(InOrderQueue); auto Node1Graph = InOrderQueue.submit( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); auto PtrNode1 = sycl::detail::getSyclObjImpl(InOrderGraph) @@ -995,7 +911,7 @@ TEST_F(CommandGraphTest, InOrderQueueWithEmpty) { ASSERT_EQ(PtrNode2->MPredecessors.front().lock(), PtrNode1); auto Node3Graph = InOrderQueue.submit( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); auto PtrNode3 = sycl::detail::getSyclObjImpl(InOrderGraph) @@ -1039,7 +955,7 @@ TEST_F(CommandGraphTest, InOrderQueueWithEmptyFirst) { ASSERT_TRUE(PtrNode1->MPredecessors.empty()); auto Node2Graph = InOrderQueue.submit( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); auto PtrNode2 = sycl::detail::getSyclObjImpl(InOrderGraph) @@ -1052,7 +968,7 @@ TEST_F(CommandGraphTest, InOrderQueueWithEmptyFirst) { ASSERT_EQ(PtrNode2->MPredecessors.front().lock(), PtrNode1); auto Node3Graph = InOrderQueue.submit( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); auto PtrNode3 = sycl::detail::getSyclObjImpl(InOrderGraph) @@ -1088,7 +1004,7 @@ TEST_F(CommandGraphTest, InOrderQueueWithEmptyLast) { // Record in-order queue with two regular nodes then an empty node InOrderGraph.begin_recording(InOrderQueue); auto Node1Graph = InOrderQueue.submit( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); auto PtrNode1 = sycl::detail::getSyclObjImpl(InOrderGraph) @@ -1097,7 +1013,7 @@ TEST_F(CommandGraphTest, InOrderQueueWithEmptyLast) { ASSERT_TRUE(PtrNode1->MPredecessors.empty()); auto Node2Graph = InOrderQueue.submit( - [&](sycl::handler &cgh) { cgh.single_task([]() {}); }); + [&](sycl::handler &cgh) { cgh.single_task>([]() {}); }); auto PtrNode2 = sycl::detail::getSyclObjImpl(InOrderGraph) @@ -1285,7 +1201,7 @@ TEST_F(CommandGraphTest, Reductions) { try { Graph.add([&](handler &CGH) { CGH.parallel_for( - range<1>{1}, reduction(&ReduVar, int{0}, sycl::plus()), + range<1>{1}, reduction(&ReduVar, int{0}, sycl::plus<>()), [=](item<1> idx, auto &Sum) {}); }); } catch (const sycl::exception &e) { @@ -1404,265 +1320,3 @@ TEST_F(CommandGraphTest, GetProfilingInfoExceptionCheck) { } ASSERT_EQ(ExceptionCode, sycl::errc::invalid); } - -class MultiThreadGraphTest : public CommandGraphTest { -public: - MultiThreadGraphTest() - : CommandGraphTest(), NumThreads(std::thread::hardware_concurrency()), - SyncPoint(NumThreads) { - Threads.reserve(NumThreads); - } - -protected: - const unsigned NumThreads; - Barrier SyncPoint; - std::vector Threads; -}; - -TEST_F(MultiThreadGraphTest, BeginEndRecording) { - auto RecordGraph = [&]() { - queue MyQueue{Queue.get_context(), Queue.get_device()}; - - SyncPoint.wait(); - - Graph.begin_recording(MyQueue); - runKernels(MyQueue); - Graph.end_recording(MyQueue); - }; - - for (unsigned i = 0; i < NumThreads; ++i) { - Threads.emplace_back(RecordGraph); - } - - for (unsigned i = 0; i < NumThreads; ++i) { - Threads[i].join(); - } - - // Reference computation - queue QueueRef; - experimental::command_graph GraphRef{ - Queue.get_context(), Queue.get_device()}; - - for (unsigned i = 0; i < NumThreads; ++i) { - queue MyQueue; - GraphRef.begin_recording(MyQueue); - runKernels(MyQueue); - GraphRef.end_recording(MyQueue); - } - - auto GraphImpl = sycl::detail::getSyclObjImpl(Graph); - auto GraphRefImpl = sycl::detail::getSyclObjImpl(GraphRef); - ASSERT_EQ(GraphImpl->hasSimilarStructure(GraphRefImpl), true); -} - -TEST_F(MultiThreadGraphTest, ExplicitAddNodes) { - auto RecordGraph = [&]() { - queue MyQueue{Queue.get_context(), Queue.get_device()}; - - SyncPoint.wait(); - addKernels(Graph); - }; - - for (unsigned i = 0; i < NumThreads; ++i) { - Threads.emplace_back(RecordGraph); - } - - for (unsigned i = 0; i < NumThreads; ++i) { - Threads[i].join(); - } - - // Reference computation - queue QueueRef; - experimental::command_graph GraphRef{ - Queue.get_context(), Queue.get_device()}; - - for (unsigned i = 0; i < NumThreads; ++i) { - addKernels(GraphRef); - } - - auto GraphImpl = sycl::detail::getSyclObjImpl(Graph); - auto GraphRefImpl = sycl::detail::getSyclObjImpl(GraphRef); - ASSERT_EQ(GraphImpl->hasSimilarStructure(GraphRefImpl), true); -} - -TEST_F(MultiThreadGraphTest, RecordAddNodes) { - Graph.begin_recording(Queue); - auto RecordGraph = [&]() { - SyncPoint.wait(); - runKernels(Queue); - }; - - for (unsigned i = 0; i < NumThreads; ++i) { - Threads.emplace_back(RecordGraph); - } - - for (unsigned i = 0; i < NumThreads; ++i) { - Threads[i].join(); - } - - // We stop recording the Queue when all threads have finished their processing - Graph.end_recording(Queue); - - // Reference computation - queue QueueRef; - experimental::command_graph GraphRef{ - Queue.get_context(), Queue.get_device()}; - - GraphRef.begin_recording(QueueRef); - for (unsigned i = 0; i < NumThreads; ++i) { - runKernels(QueueRef); - } - GraphRef.end_recording(QueueRef); - - auto GraphImpl = sycl::detail::getSyclObjImpl(Graph); - auto GraphRefImpl = sycl::detail::getSyclObjImpl(GraphRef); - ASSERT_EQ(GraphImpl->hasSimilarStructure(GraphRefImpl), true); -} - -TEST_F(MultiThreadGraphTest, RecordAddNodesInOrderQueue) { - sycl::property_list Properties{sycl::property::queue::in_order()}; - queue InOrderQueue{Dev, Properties}; - - experimental::command_graph - InOrderGraph{InOrderQueue.get_context(), InOrderQueue.get_device()}; - - InOrderGraph.begin_recording(InOrderQueue); - auto RecordGraph = [&]() { - SyncPoint.wait(); - runKernelsInOrder(InOrderQueue); - }; - - for (unsigned i = 0; i < NumThreads; ++i) { - Threads.emplace_back(RecordGraph); - } - - for (unsigned i = 0; i < NumThreads; ++i) { - Threads[i].join(); - } - - // We stop recording the Queue when all threads have finished their processing - InOrderGraph.end_recording(InOrderQueue); - - // Reference computation - queue InOrderQueueRef{Dev, Properties}; - experimental::command_graph - InOrderGraphRef{InOrderQueueRef.get_context(), - InOrderQueueRef.get_device()}; - - InOrderGraphRef.begin_recording(InOrderQueueRef); - for (unsigned i = 0; i < NumThreads; ++i) { - runKernelsInOrder(InOrderQueueRef); - } - InOrderGraphRef.end_recording(InOrderQueueRef); - - auto GraphImpl = sycl::detail::getSyclObjImpl(InOrderGraph); - auto GraphRefImpl = sycl::detail::getSyclObjImpl(InOrderGraphRef); - ASSERT_EQ(GraphImpl->getNumberOfNodes(), GraphRefImpl->getNumberOfNodes()); - - // In-order graph must have only a single root - ASSERT_EQ(GraphImpl->MRoots.size(), 1lu); - - // Check structure graph - for (auto Node : GraphImpl->MRoots) { - ASSERT_EQ(depthSearchSuccessorCheck(Node), true); - } -} - -TEST_F(MultiThreadGraphTest, Finalize) { - addKernels(Graph); - - std::map> - GraphsExecMap; - auto FinalizeGraph = [&](int ThreadNum) { - SyncPoint.wait(); - auto GraphExec = Graph.finalize(); - - GraphsExecMap.insert( - std::map>:: - value_type(ThreadNum, GraphExec)); - Queue.submit([&](sycl::handler &CGH) { CGH.ext_oneapi_graph(GraphExec); }); - }; - - for (unsigned i = 0; i < NumThreads; ++i) { - Threads.emplace_back(FinalizeGraph, i); - } - - for (unsigned i = 0; i < NumThreads; ++i) { - Threads[i].join(); - } - - // Reference computation - queue QueueRef; - experimental::command_graph GraphRef{ - Queue.get_context(), Queue.get_device()}; - - addKernels(GraphRef); - - for (unsigned i = 0; i < NumThreads; ++i) { - auto GraphExecRef = GraphRef.finalize(); - QueueRef.submit( - [&](sycl::handler &CGH) { CGH.ext_oneapi_graph(GraphExecRef); }); - auto GraphExecImpl = - sycl::detail::getSyclObjImpl(GraphsExecMap.find(i)->second); - auto GraphExecRefImpl = sycl::detail::getSyclObjImpl(GraphExecRef); - ASSERT_EQ(checkExecGraphSchedule(GraphExecImpl, GraphExecRefImpl), true); - } -} - -TEST_F(CommandGraphTest, InvalidBuffer) { - // Check that using a buffer with write_back enabled in a graph will throw. - int Data; - // Create a buffer which does not have write-back disabled. - buffer Buffer{&Data, range<1>{1}}; - - // Use this buffer in the graph, this should throw. - ASSERT_THROW( - { - try { - Graph.add([&](handler &CGH) { - auto Acc = Buffer.get_access(CGH); - }); - } catch (const sycl::exception &e) { - ASSERT_EQ(e.code(), make_error_code(sycl::errc::invalid)); - throw; - } - }, - sycl::exception); -} - -TEST_F(CommandGraphTest, InvalidHostAccessor) { - // Check that creating a host_accessor on a buffer which is in use by a graph - // will throw. - - // Create a buffer which does not have write-back disabled. - buffer Buffer{range<1>{1}}; - - { - // Create a graph in local scope so we can destroy it - ext::oneapi::experimental::command_graph Graph{ - Queue.get_context(), - Queue.get_device(), - {experimental::property::graph::assume_buffer_outlives_graph{}}}; - - // Add the buffer to the graph. - Graph.add([&](handler &CGH) { - auto Acc = Buffer.get_access(CGH); - }); - - // Attempt to create a host_accessor, which should throw. - ASSERT_THROW( - { - try { - host_accessor HostAcc{Buffer}; - } catch (const sycl::exception &e) { - ASSERT_EQ(e.code(), make_error_code(sycl::errc::invalid)); - throw; - } - }, - sycl::exception); - } - // Graph is now out of scope so we should be able to create a host_accessor - ASSERT_NO_THROW({ host_accessor HostAcc{Buffer}; }); -} diff --git a/sycl/unittests/helpers/PiMockPlugin.hpp b/sycl/unittests/helpers/PiMockPlugin.hpp index 5c1241ca0d49e..c393d57d16d6c 100644 --- a/sycl/unittests/helpers/PiMockPlugin.hpp +++ b/sycl/unittests/helpers/PiMockPlugin.hpp @@ -164,7 +164,7 @@ inline pi_result mock_piDeviceGetInfo(pi_device device, size_t *param_value_size_ret) { constexpr char MockDeviceName[] = "Mock device"; constexpr char MockSupportedExtensions[] = - "cl_khr_fp64 cl_khr_fp16 cl_khr_il_program"; + "cl_khr_fp64 cl_khr_fp16 cl_khr_il_program ur_exp_command_buffer"; switch (param_name) { case PI_DEVICE_INFO_TYPE: { // Act like any device is a GPU. From 6dd800738e31460e374155d6b9f62a702a44d570 Mon Sep 17 00:00:00 2001 From: Ewan Crawford Date: Wed, 16 Aug 2023 12:03:56 +0100 Subject: [PATCH 07/10] Give enum class entries a value --- .../sycl/ext/oneapi/experimental/graph.hpp | 16 ++++++++-------- 1 file changed, 8 insertions(+), 8 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/graph.hpp b/sycl/include/sycl/ext/oneapi/experimental/graph.hpp index 40f7c4cbbe51d..de41a6b28b781 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/graph.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/graph.hpp @@ -34,14 +34,14 @@ namespace detail { // List of sycl features and extensions which are not supported by graphs. Used // for throwing errors when these features are used with graphs. enum class UnsupportedGraphFeatures { - sycl_reductions, - sycl_specialization_constants, - sycl_kernel_bundle, - sycl_ext_oneapi_kernel_properties, - sycl_ext_oneapi_enqueue_barrier, - sycl_ext_oneapi_memcpy2d, - sycl_ext_oneapi_device_global, - sycl_ext_oneapi_bindless_images + sycl_reductions = 0, + sycl_specialization_constants = 1, + sycl_kernel_bundle = 2, + sycl_ext_oneapi_kernel_properties = 3, + sycl_ext_oneapi_enqueue_barrier = 4, + sycl_ext_oneapi_memcpy2d = 5, + sycl_ext_oneapi_device_global = 6, + sycl_ext_oneapi_bindless_images = 7 }; constexpr const char * From e7466639c02ab2f217d0cef22ff294ce064d67c9 Mon Sep 17 00:00:00 2001 From: Ewan Crawford Date: Thu, 17 Aug 2023 17:08:55 +0100 Subject: [PATCH 08/10] Use M prefix for new member variable --- sycl/source/detail/event_impl.cpp | 2 +- sycl/source/detail/event_impl.hpp | 6 +++--- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/sycl/source/detail/event_impl.cpp b/sycl/source/detail/event_impl.cpp index 21a1e720dd070..f3a5eabe04976 100644 --- a/sycl/source/detail/event_impl.cpp +++ b/sycl/source/detail/event_impl.cpp @@ -274,7 +274,7 @@ void event_impl::checkProfilingPreconditions() const { "Profiling information is unavailable as the queue associated with " "the event does not have the 'enable_profiling' property."); } - if (EventFromSubmitedExecCommandBuffer) { + if (MEventFromSubmitedExecCommandBuffer) { throw sycl::exception(make_error_code(sycl::errc::invalid), "Profiling information is unavailable for events " "returned by a graph submission."); diff --git a/sycl/source/detail/event_impl.hpp b/sycl/source/detail/event_impl.hpp index 2bcad45bd9d33..31ffd23bba01c 100644 --- a/sycl/source/detail/event_impl.hpp +++ b/sycl/source/detail/event_impl.hpp @@ -279,11 +279,11 @@ class event_impl { } void setEventFromSubmitedExecCommandBuffer(bool value) { - EventFromSubmitedExecCommandBuffer = value; + MEventFromSubmitedExecCommandBuffer = value; } bool isEventFromSubmitedExecCommandBuffer() const { - return EventFromSubmitedExecCommandBuffer; + return MEventFromSubmitedExecCommandBuffer; } protected: @@ -336,7 +336,7 @@ class event_impl { /// This event is also be stored in the graph so a weak_ptr is used. std::weak_ptr MGraph; /// Indicates that the event results from a command graph submission - bool EventFromSubmitedExecCommandBuffer = false; + bool MEventFromSubmitedExecCommandBuffer = false; // If this event represents a submission to a // sycl::detail::pi::PiExtCommandBuffer the sync point for that submission is From 593e5f5916c9d07a19b2cabdf5cd4c73e917223b Mon Sep 17 00:00:00 2001 From: Ewan Crawford Date: Mon, 21 Aug 2023 12:09:00 +0100 Subject: [PATCH 09/10] Fix CI warning in reduction.hpp Using a sycl::reduction in a graphs unittests caused the build to fail with a sign compare warning ``` /__w/llvm/llvm/build/include/sycl/reduction.hpp:1358:56: error: comparison of integer expressions of different signedness: 'int' and 'const size_t' {aka 'const long unsigned int'} [-Werror=sign-compare] ``` --- sycl/include/sycl/reduction.hpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/sycl/include/sycl/reduction.hpp b/sycl/include/sycl/reduction.hpp index 77c60b9bb7829..c29e109128f35 100644 --- a/sycl/include/sycl/reduction.hpp +++ b/sycl/include/sycl/reduction.hpp @@ -1355,7 +1355,8 @@ struct NDRangeReduction< sycl::atomic_ref( NWorkGroupsFinished[0]); - DoReducePartialSumsInLastWG[0] = ++NFinished == NWorkGroups; + DoReducePartialSumsInLastWG[0] = + ++NFinished == static_cast(NWorkGroups); } workGroupBarrier(); From 163607f17de9d45071e2a0abf9cd661ffd4ae92a Mon Sep 17 00:00:00 2001 From: Ewan Crawford Date: Mon, 21 Aug 2023 13:27:35 +0100 Subject: [PATCH 10/10] Remove symbol from linux ABI test --- sycl/test/abi/sycl_symbols_linux.dump | 1 - 1 file changed, 1 deletion(-) diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 290a08bf31e14..fcdd5aa8171c0 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -4685,7 +4685,6 @@ _ZNK4sycl3_V17context8get_infoINS0_4info7context8platformEEENS0_6detail20is_cont _ZNK4sycl3_V17context9getNativeEv _ZNK4sycl3_V17handler15getCommandGraphEv _ZNK4sycl3_V17handler17getContextImplPtrEv -_ZNK4sycl3_V17handler22throwIfGraphAssociatedILNS0_3ext6oneapi12experimental6detail24UnsupportedGraphFeaturesE7EEEvv _ZNK4sycl3_V17handler27isStateExplicitKernelBundleEv _ZNK4sycl3_V17handler30getOrInsertHandlerKernelBundleEb _ZNK4sycl3_V17sampler12get_propertyINS0_3ext6oneapi4cuda8property7context19use_primary_contextEEET_v