From d1a6befceb022f7e1e8c58ffeb4c6ce72ccd97d2 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?F=C3=A1bio=20Mestre?= Date: Wed, 8 Jan 2025 14:00:06 +0000 Subject: [PATCH 01/10] Implement Dynamic Local Accessors --- .../sycl/ext/oneapi/experimental/graph.hpp | 48 +++++++++ sycl/include/sycl/handler.hpp | 23 ++++ sycl/source/detail/graph_impl.cpp | 88 +++++++++++++++ sycl/source/detail/graph_impl.hpp | 40 +++++++ .../Update/dyn_cgf_dyn_local_accessor.cpp | 85 +++++++++++++++ .../Update/update_dynamic_local_accessor.cpp | 71 ++++++++++++ ...amic_local_accessor_multiple_accessors.cpp | 90 ++++++++++++++++ ..._dynamic_local_accessor_multiple_nodes.cpp | 102 ++++++++++++++++++ .../Extensions/CommandGraph/Update.cpp | 37 +++++++ 9 files changed, 584 insertions(+) create mode 100644 sycl/test-e2e/Graph/Update/dyn_cgf_dyn_local_accessor.cpp create mode 100644 sycl/test-e2e/Graph/Update/update_dynamic_local_accessor.cpp create mode 100644 sycl/test-e2e/Graph/Update/update_dynamic_local_accessor_multiple_accessors.cpp create mode 100644 sycl/test-e2e/Graph/Update/update_dynamic_local_accessor_multiple_nodes.cpp diff --git a/sycl/include/sycl/ext/oneapi/experimental/graph.hpp b/sycl/include/sycl/ext/oneapi/experimental/graph.hpp index e2e87c30ea945..898e0d06e595e 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/graph.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/graph.hpp @@ -447,6 +447,11 @@ class command_graph namespace detail { class __SYCL_EXPORT dynamic_parameter_base { public: + + dynamic_parameter_base( + sycl::ext::oneapi::experimental::command_graph + Graph); + dynamic_parameter_base( sycl::ext::oneapi::experimental::command_graph Graph, @@ -461,6 +466,13 @@ class __SYCL_EXPORT dynamic_parameter_base { void updateValue(const raw_kernel_arg *NewRawValue, size_t Size); void updateAccessor(const sycl::detail::AccessorBaseHost *Acc); + + sycl::detail::LocalAccessorImplPtr getLocalAccessor(handler* Handler); + + void registerLocalAccessor(sycl::detail::LocalAccessorBaseHost* LocalAccBaseHost, handler* Handler); + + void updateLocalAccessor(range<3> NewAllocationSize); + std::shared_ptr impl; template @@ -498,6 +510,42 @@ class dynamic_parameter : public detail::dynamic_parameter_base { } }; +template +class dynamic_local_accessor : public detail::dynamic_parameter_base { +public: + template 0)>> + dynamic_local_accessor(command_graph Graph, + range AllocationSize, + const property_list &PropList = {}) + : detail::dynamic_parameter_base(Graph), AllocationSize(AllocationSize) { + (void)PropList; + } + + void update(range NewAllocationSize) { + detail::dynamic_parameter_base::updateLocalAccessor( + ::sycl::detail::convertToArrayOfN<3, 1>(NewAllocationSize)); + }; + + local_accessor get(handler &CGH) { +#ifndef __SYCL_DEVICE_ONLY__ + ::sycl::detail::LocalAccessorImplPtr BaseLocalAcc = getLocalAccessor(&CGH); + if (BaseLocalAcc) { + return sycl::detail::createSyclObjFromImpl>(BaseLocalAcc); + } else { + local_accessor LocalAccessor(AllocationSize, CGH); + registerLocalAccessor( + static_cast(&LocalAccessor), &CGH); + return LocalAccessor; + } +#else + return local_accessor(); +#endif + }; + +private: + range AllocationSize; +}; + /// Additional CTAD deduction guides. template dynamic_parameter(experimental::command_graph Graph, diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 758daa3a81a9b..4b6cc89c14ed9 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -647,6 +647,22 @@ class __SYCL_EXPORT handler { registerDynamicParameter(DynamicParam, ArgIndex); } + // setArgHelper for graph dynamic_local_accessors. + template + void + setArgHelper(int ArgIndex, + ext::oneapi::experimental::dynamic_local_accessor + &DynamicLocalAccessor) { +#ifndef __SYCL_DEVICE_ONLY__ + auto LocalAccessor = DynamicLocalAccessor.get(*this); + setArgHelper(ArgIndex, LocalAccessor); + registerDynamicParameter(DynamicLocalAccessor, ArgIndex); +#else + (void)ArgIndex; + (void)DynamicLocalAccessor; +#endif + } + // setArgHelper for the raw_kernel_arg extension type. void setArgHelper(int ArgIndex, sycl::ext::oneapi::experimental::raw_kernel_arg &&Arg) { @@ -1839,6 +1855,13 @@ class __SYCL_EXPORT handler { setArgHelper(argIndex, dynamicParam); } + template + void set_arg(int argIndex, + ext::oneapi::experimental::dynamic_local_accessor + &DynamicLocalAccessor) { + setArgHelper(argIndex, DynamicLocalAccessor); + } + // set_arg for the raw_kernel_arg extension type. void set_arg(int argIndex, ext::oneapi::experimental::raw_kernel_arg &&Arg) { setArgHelper(argIndex, std::move(Arg)); diff --git a/sycl/source/detail/graph_impl.cpp b/sycl/source/detail/graph_impl.cpp index e6181a559d8e6..9baccd6b504fb 100644 --- a/sycl/source/detail/graph_impl.cpp +++ b/sycl/source/detail/graph_impl.cpp @@ -1899,6 +1899,11 @@ dynamic_parameter_base::dynamic_parameter_base( : impl(std::make_shared( sycl::detail::getSyclObjImpl(Graph), ParamSize, Data)) {} +dynamic_parameter_base::dynamic_parameter_base( + command_graph Graph) + : impl(std::make_shared( + sycl::detail::getSyclObjImpl(Graph))) {} + void dynamic_parameter_base::updateValue(const void *NewValue, size_t Size) { impl->updateValue(NewValue, Size); } @@ -1913,6 +1918,20 @@ void dynamic_parameter_base::updateAccessor( impl->updateAccessor(Acc); } +sycl::detail::LocalAccessorImplPtr +dynamic_parameter_base::getLocalAccessor(handler *Handler) { + return impl->getLocalAccessor(Handler); +} + +void dynamic_parameter_base::registerLocalAccessor( + sycl::detail::LocalAccessorBaseHost *LocalAccBaseHost, handler *Handler) { + impl->registerLocalAccessor(LocalAccBaseHost, Handler); +} + +void dynamic_parameter_base::updateLocalAccessor(range<3> NewAllocationSize) { + impl->updateLocalAccessor(NewAllocationSize); +} + void dynamic_parameter_impl::updateValue(const raw_kernel_arg *NewRawValue, size_t Size) { // Number of bytes is taken from member of raw_kernel_arg object rather @@ -1968,6 +1987,53 @@ void dynamic_parameter_impl::updateAccessor( sizeof(sycl::detail::AccessorBaseHost)); } +sycl::detail::LocalAccessorImplPtr +dynamic_parameter_impl::getLocalAccessor(handler *Handler) { + auto HandlerImpl = sycl::detail::getSyclObjImpl(*Handler); + auto FindLocalAcc = MHandlerToLocalAccMap.find(HandlerImpl); + + if (FindLocalAcc != MHandlerToLocalAccMap.end()) { + auto LocalAccImpl = FindLocalAcc->second; + return LocalAccImpl; + } + return nullptr; +} + +void dynamic_parameter_impl::registerLocalAccessor( + sycl::detail::LocalAccessorBaseHost *LocalAccBaseHost, handler *Handler) { + + auto HandlerImpl = sycl::detail::getSyclObjImpl(*Handler); + auto LocalAccImpl = sycl::detail::getSyclObjImpl(*LocalAccBaseHost); + + MHandlerToLocalAccMap.insert({HandlerImpl, LocalAccImpl}); +} + +void dynamic_parameter_impl::updateLocalAccessor(range<3> NewAllocationSize) { + + for (auto &[NodeWeak, ArgIndex] : MNodes) { + auto NodeShared = NodeWeak.lock(); + if (NodeShared) { + // We can use the first local accessor in the map since the dimensions + // and element type should be identical. + auto LocalAccessor = MHandlerToLocalAccMap.begin()->second; + dynamic_parameter_impl::updateCGLocalAccessor( + NodeShared->MCommandGroup, ArgIndex, NewAllocationSize, + LocalAccessor->MDims, LocalAccessor->MElemSize); + } + } + + for (auto &DynCGInfo : MDynCGs) { + auto DynCG = DynCGInfo.DynCG.lock(); + if (DynCG) { + auto &CG = DynCG->MKernels[DynCGInfo.CGIndex]; + auto LocalAccessor = MHandlerToLocalAccMap.begin()->second; + dynamic_parameter_impl::updateCGLocalAccessor( + CG, DynCGInfo.ArgIndex, NewAllocationSize, LocalAccessor->MDims, + LocalAccessor->MElemSize); + } + } +} + void dynamic_parameter_impl::updateCGArgValue( std::shared_ptr CG, int ArgIndex, const void *NewValue, size_t Size) { @@ -2033,6 +2099,27 @@ void dynamic_parameter_impl::updateCGAccessor( } } +void dynamic_parameter_impl::updateCGLocalAccessor( + std::shared_ptr CG, int ArgIndex, + range<3> NewAllocationSize, int Dims, int ElemSize) { + auto &Args = static_cast(CG.get())->MArgs; + + for (auto &Arg : Args) { + if (Arg.MIndex != ArgIndex) { + continue; + } + assert(Arg.MType == sycl::detail::kernel_param_kind_t::kind_std_layout); + + int SizeInBytes = ElemSize; + for (int I = 0; I < Dims; ++I) + SizeInBytes *= NewAllocationSize[I]; + SizeInBytes = std::max(SizeInBytes, 1); + + Arg.MSize = SizeInBytes; + break; + } +} + dynamic_command_group_impl::dynamic_command_group_impl( const command_graph &Graph) : MGraph{sycl::detail::getSyclObjImpl(Graph)}, MActiveCGF(0) {} @@ -2154,6 +2241,7 @@ size_t dynamic_command_group::get_active_index() const { void dynamic_command_group::set_active_index(size_t Index) { return impl->setActiveIndex(Index); } + } // namespace experimental } // namespace oneapi } // namespace ext diff --git a/sycl/source/detail/graph_impl.hpp b/sycl/source/detail/graph_impl.hpp index e609123b4f285..993e1644a22fd 100644 --- a/sycl/source/detail/graph_impl.hpp +++ b/sycl/source/detail/graph_impl.hpp @@ -1412,6 +1412,10 @@ class exec_graph_impl { class dynamic_parameter_impl { public: + /// Used for parameters that don't have data such as local_accessors. + dynamic_parameter_impl(std::shared_ptr GraphImpl) + : MGraph(GraphImpl) {} + dynamic_parameter_impl(std::shared_ptr GraphImpl, size_t ParamSize, const void *Data) : MGraph(GraphImpl), MValueStorage(ParamSize) { @@ -1477,6 +1481,26 @@ class dynamic_parameter_impl { /// @param Acc The new accessor value void updateAccessor(const sycl::detail::AccessorBaseHost *Acc); + /// Updates the value of all local accessors in registered nodes and dynamic + /// CGs. + /// @param NewAllocationSize The new size for the update local accessors. + void updateLocalAccessor(range<3> NewAllocationSize); + + /// Gets the implementation for the local accessor that is associated with + /// a specific handler. + /// @param The handler that the local accessor is associated with. + /// @return returns the impl object for the local accessor that is associated + /// with this handler. Or nullptr if no local accessor has been registered + /// for this handler. + sycl::detail::LocalAccessorImplPtr getLocalAccessor(handler *Handler); + + /// Associates a local accessor with this dynamic local accessor for a + /// specific handler. + /// @param LocalAccBase the local accessor that needs to be registered. + /// @param Handler the handler that the LocalAccessor is associated with. + void registerLocalAccessor(sycl::detail::LocalAccessorBaseHost *LocalAccBase, + handler *Handler); + /// Static helper function for updating command-group value arguments. /// @param CG The command-group to update the argument information for. /// @param ArgIndex The argument index to update. @@ -1493,6 +1517,18 @@ class dynamic_parameter_impl { int ArgIndex, const sycl::detail::AccessorBaseHost *Acc); + /// Static helper function for updating command-group local accessor + /// arguments. + /// @param CG The command-group to update the argument information for. + /// @param ArgIndex The argument index to update. + /// @param NewAllocationSize The new allocation size for the local accessor + /// argument. + /// @param Dims The dimensions of the local accessor argument. + /// @param ElemSize The size of each element in the local accessor. + static void updateCGLocalAccessor(std::shared_ptr CG, + int ArgIndex, range<3> NewAllocationSize, + int Dims, int ElemSize); + // Weak ptrs to node_impls which will be updated std::vector, int>> MNodes; // Dynamic command-groups which will be updated @@ -1500,6 +1536,10 @@ class dynamic_parameter_impl { std::shared_ptr MGraph; std::vector MValueStorage; + + std::unordered_map, + sycl::detail::LocalAccessorImplPtr> + MHandlerToLocalAccMap; }; class dynamic_command_group_impl diff --git a/sycl/test-e2e/Graph/Update/dyn_cgf_dyn_local_accessor.cpp b/sycl/test-e2e/Graph/Update/dyn_cgf_dyn_local_accessor.cpp new file mode 100644 index 0000000000000..91cd54e7fb14c --- /dev/null +++ b/sycl/test-e2e/Graph/Update/dyn_cgf_dyn_local_accessor.cpp @@ -0,0 +1,85 @@ +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out +// Extra run to check for leaks in Level Zero using UR_L0_LEAKS_DEBUG +// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=0 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %} +// Extra run to check for immediate-command-list in Level Zero +// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %} + +// Tests using dynamic command-group objects with dynamic local accessors. + +#include "../graph_common.hpp" + +int main() { + using T = int; + + const size_t LocalMemSize = 128; + + queue Queue{}; + + std::vector HostDataBeforeUpdate(Size); + std::vector HostDataAfterUpdate(Size); + std::iota(HostDataBeforeUpdate.begin(), HostDataBeforeUpdate.end(), 10); + + T *PtrA = malloc_device(Size, Queue); + Queue.copy(HostDataBeforeUpdate.data(), PtrA, Size); + Queue.wait_and_throw(); + + exp_ext::command_graph Graph{Queue.get_context(), Queue.get_device()}; + + exp_ext::dynamic_local_accessor DynLocalAccessor{Graph, LocalMemSize}; + + auto CGFA = [&](handler &CGH) { + CGH.set_arg(0, DynLocalAccessor); + auto LocalMem = DynLocalAccessor.get(CGH); + + CGH.parallel_for(nd_range({Size}, {LocalMemSize}), [=](nd_item<1> Item) { + LocalMem[Item.get_local_linear_id()] = Item.get_local_linear_id(); + PtrA[Item.get_global_linear_id()] = LocalMem[Item.get_local_linear_id()]; + }); + }; + + auto CGFB = [&](handler &CGH) { + CGH.set_arg(0, DynLocalAccessor); + auto LocalMem = DynLocalAccessor.get(CGH); + + CGH.parallel_for( + nd_range({Size}, {LocalMemSize * 2}), [=](nd_item<1> Item) { + LocalMem[Item.get_local_linear_id()] = Item.get_local_linear_id(); + PtrA[Item.get_global_linear_id()] = + LocalMem[Item.get_local_linear_id()] * 2; + }); + }; + + auto DynamicCG = exp_ext::dynamic_command_group(Graph, {CGFA, CGFB}); + auto DynamicCGNode = Graph.add(DynamicCG); + auto ExecGraph = Graph.finalize(exp_ext::property::graph::updatable{}); + + Queue.submit([&](handler &CGH) { CGH.ext_oneapi_graph(ExecGraph); }); + Queue.wait_and_throw(); + Queue.copy(PtrA, HostDataBeforeUpdate.data(), Size); + Queue.wait_and_throw(); + + DynLocalAccessor.update(LocalMemSize * 2); + DynamicCG.set_active_index(1); + ExecGraph.update(DynamicCGNode); + Queue.ext_oneapi_graph(ExecGraph).wait(); + + Queue.submit([&](handler &CGH) { CGH.ext_oneapi_graph(ExecGraph); }); + Queue.wait_and_throw(); + Queue.copy(PtrA, HostDataAfterUpdate.data(), Size); + Queue.wait_and_throw(); + + for (size_t i = 0; i < Size; i++) { + T Ref = i % LocalMemSize; + assert(check_value(i, Ref, HostDataBeforeUpdate[i], "PtrA Before Update")); + } + + for (size_t i = 0; i < Size; i++) { + T Ref = i % (LocalMemSize * 2) * 2; + assert(check_value(i, Ref, HostDataAfterUpdate[i], "PtrA After Update")); + } + + free(PtrA, Queue); + + return 0; +} diff --git a/sycl/test-e2e/Graph/Update/update_dynamic_local_accessor.cpp b/sycl/test-e2e/Graph/Update/update_dynamic_local_accessor.cpp new file mode 100644 index 0000000000000..2b485e53d9783 --- /dev/null +++ b/sycl/test-e2e/Graph/Update/update_dynamic_local_accessor.cpp @@ -0,0 +1,71 @@ +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out +// Extra run to check for leaks in Level Zero using UR_L0_LEAKS_DEBUG +// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=0 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %} +// Extra run to check for immediate-command-list in Level Zero +// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %} + +// Tests updating local accessor parameters. +#include "../graph_common.hpp" + +int main() { + using T = int; + + const size_t LocalMemSize = 128; + + queue Queue{}; + + std::vector HostDataBeforeUpdate(Size); + std::vector HostDataAfterUpdate(Size); + std::iota(HostDataBeforeUpdate.begin(), HostDataBeforeUpdate.end(), 10); + + T *PtrA = malloc_device(Size, Queue); + Queue.copy(HostDataBeforeUpdate.data(), PtrA, Size); + Queue.wait_and_throw(); + + exp_ext::command_graph Graph{Queue.get_context(), Queue.get_device()}; + + exp_ext::dynamic_local_accessor DynLocalAccessor{Graph, LocalMemSize}; + + auto Node = Graph.add([&](handler &CGH) { + CGH.set_arg(0, DynLocalAccessor); + auto LocalMem = DynLocalAccessor.get(CGH); + + CGH.parallel_for(nd_range({Size}, {LocalMemSize}), [=](nd_item<1> Item) { + LocalMem[Item.get_local_linear_id()] = Item.get_local_linear_id(); + PtrA[Item.get_global_linear_id()] = LocalMem[Item.get_local_linear_id()]; + }); + }); + + auto GraphExec = Graph.finalize(exp_ext::property::graph::updatable{}); + + // Submit the graph before the update and save the results. + Queue.submit([&](handler &CGH) { CGH.ext_oneapi_graph(GraphExec); }); + Queue.wait_and_throw(); + Queue.copy(PtrA, HostDataBeforeUpdate.data(), Size); + Queue.wait_and_throw(); + + DynLocalAccessor.update(LocalMemSize * 2); + Node.update_nd_range(nd_range({Size}, {LocalMemSize * 2})); + GraphExec.update(Node); + + // Submit the graph after the update and save the results. + Queue.submit([&](handler &CGH) { CGH.ext_oneapi_graph(GraphExec); }); + Queue.wait_and_throw(); + Queue.copy(PtrA, HostDataAfterUpdate.data(), Size); + Queue.wait_and_throw(); + + for (size_t i = 0; i < Size; i++) { + T Ref = i % LocalMemSize; + assert(check_value(i, Ref, HostDataBeforeUpdate[i], "PtrA Before Update")); + } + + for (size_t i = 0; i < Size; i++) { + T Ref = i % (LocalMemSize * 2); + assert(check_value(i, Ref, HostDataAfterUpdate[i], "PtrA After Update")); + } + + free(PtrA, Queue); + + return 0; +} diff --git a/sycl/test-e2e/Graph/Update/update_dynamic_local_accessor_multiple_accessors.cpp b/sycl/test-e2e/Graph/Update/update_dynamic_local_accessor_multiple_accessors.cpp new file mode 100644 index 0000000000000..d3bf84c5f5f64 --- /dev/null +++ b/sycl/test-e2e/Graph/Update/update_dynamic_local_accessor_multiple_accessors.cpp @@ -0,0 +1,90 @@ +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out +// Extra run to check for leaks in Level Zero using UR_L0_LEAKS_DEBUG +// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=0 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %} +// Extra run to check for immediate-command-list in Level Zero +// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %} + +// Tests updating local accessor parameters using multiple dynamic local +// accessors in the graph node. +#include "../graph_common.hpp" + +int main() { + using T = int; + + const size_t LocalMemSize = 128; + + queue Queue{}; + + std::vector HostDataBeforeUpdate(Size); + std::vector HostDataAfterUpdate(Size); + std::iota(HostDataBeforeUpdate.begin(), HostDataBeforeUpdate.end(), 10); + + T *PtrA = malloc_device(Size, Queue); + Queue.copy(HostDataBeforeUpdate.data(), PtrA, Size); + Queue.wait_and_throw(); + + exp_ext::command_graph Graph{Queue.get_context(), Queue.get_device()}; + + exp_ext::dynamic_local_accessor DynLocalAccessor{Graph, LocalMemSize}; + exp_ext::dynamic_local_accessor DynLocalAccessor2{Graph, LocalMemSize}; + + auto Node = Graph.add([&](handler &CGH) { + CGH.set_arg(0, DynLocalAccessor); + CGH.set_arg(5, DynLocalAccessor2); + CGH.set_arg(9, DynLocalAccessor); + + auto LocalMem = DynLocalAccessor.get(CGH); + + // Tests using 2 different dynamic local accessors in the same CGF. + auto LocalMem2 = DynLocalAccessor2.get(CGH); + + // Tests getting another local accessor from a dynamic local accessor that + // was already used in this CGF. + auto LocalMem3 = DynLocalAccessor.get(CGH); + + CGH.parallel_for(nd_range({Size}, {LocalMemSize}), [=](nd_item<1> Item) { + LocalMem[Item.get_local_linear_id()] = Item.get_local_linear_id(); + PtrA[Item.get_global_linear_id()] = LocalMem[Item.get_local_linear_id()]; + LocalMem2[Item.get_local_linear_id()] = Item.get_local_linear_id(); + PtrA[Item.get_global_linear_id()] += + LocalMem2[Item.get_local_linear_id()]; + LocalMem3[Item.get_local_linear_id()] = Item.get_local_linear_id(); + PtrA[Item.get_global_linear_id()] += + LocalMem3[Item.get_local_linear_id()]; + }); + }); + + auto GraphExec = Graph.finalize(exp_ext::property::graph::updatable{}); + + // Submit the graph before the update and save the results. + Queue.submit([&](handler &CGH) { CGH.ext_oneapi_graph(GraphExec); }); + Queue.wait_and_throw(); + Queue.copy(PtrA, HostDataBeforeUpdate.data(), Size); + Queue.wait_and_throw(); + + DynLocalAccessor.update(LocalMemSize * 2); + DynLocalAccessor2.update(LocalMemSize * 2); + Node.update_nd_range(nd_range({Size}, {LocalMemSize * 2})); + GraphExec.update(Node); + + // Submit the graph after the update and save the results. + Queue.submit([&](handler &CGH) { CGH.ext_oneapi_graph(GraphExec); }); + Queue.wait_and_throw(); + Queue.copy(PtrA, HostDataAfterUpdate.data(), Size); + Queue.wait_and_throw(); + + for (size_t i = 0; i < Size; i++) { + T Ref = (i % LocalMemSize) * 3; + assert(check_value(i, Ref, HostDataBeforeUpdate[i], "PtrA Before Update")); + } + + for (size_t i = 0; i < Size; i++) { + T Ref = (i % (LocalMemSize * 2)) * 3; + assert(check_value(i, Ref, HostDataAfterUpdate[i], "PtrA After Update")); + } + + free(PtrA, Queue); + + return 0; +} diff --git a/sycl/test-e2e/Graph/Update/update_dynamic_local_accessor_multiple_nodes.cpp b/sycl/test-e2e/Graph/Update/update_dynamic_local_accessor_multiple_nodes.cpp new file mode 100644 index 0000000000000..a10ae01207b2a --- /dev/null +++ b/sycl/test-e2e/Graph/Update/update_dynamic_local_accessor_multiple_nodes.cpp @@ -0,0 +1,102 @@ +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out +// Extra run to check for leaks in Level Zero using UR_L0_LEAKS_DEBUG +// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=0 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %} +// Extra run to check for immediate-command-list in Level Zero +// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %} + +// Tests updating local 2D local accessor parameters in multiple graph nodes +// simultaneously. Also tests using dynamic local accessors with +// non-zero indices. +#include "../graph_common.hpp" + +int main() { + using T = int; + + const size_t LocalMemSize = 128; + + queue Queue{}; + + std::vector HostDataBeforeUpdate(Size); + std::vector HostDataAfterUpdate(Size); + std::iota(HostDataBeforeUpdate.begin(), HostDataBeforeUpdate.end(), 10); + + T *PtrA = malloc_device(Size, Queue); + T *PtrB = malloc_device(Size, Queue); + Queue.copy(HostDataBeforeUpdate.data(), PtrA, Size); + Queue.copy(HostDataBeforeUpdate.data(), PtrB, Size); + Queue.wait_and_throw(); + + exp_ext::command_graph Graph{Queue.get_context(), Queue.get_device()}; + + exp_ext::dynamic_local_accessor DynLocalAccessor{ + Graph, range<2>{LocalMemSize, 2}}; + + auto NodeA = Graph.add([&](handler &CGH) { + CGH.set_arg(1, DynLocalAccessor); + auto LocalMem = DynLocalAccessor.get(CGH); + + CGH.parallel_for(nd_range({Size}, {LocalMemSize}), [=](nd_item<1> Item) { + PtrA[Item.get_global_linear_id()] = 0; + LocalMem[Item.get_local_linear_id()][0] = Item.get_local_linear_id(); + LocalMem[Item.get_local_linear_id()][1] = 2; + PtrA[Item.get_global_linear_id()] = + LocalMem[Item.get_local_linear_id()][0] * + LocalMem[Item.get_local_linear_id()][1]; + }); + }); + + auto NodeB = Graph.add( + [&](handler &CGH) { + CGH.set_arg(0, DynLocalAccessor); + auto LocalMem = DynLocalAccessor.get(CGH); + + CGH.parallel_for(nd_range({Size}, {LocalMemSize}), + [=](nd_item<1> Item) { + LocalMem[Item.get_local_linear_id()][0] = + Item.get_local_linear_id(); + LocalMem[Item.get_local_linear_id()][1] = 4; + PtrA[Item.get_global_linear_id()] += + LocalMem[Item.get_local_linear_id()][0] * + LocalMem[Item.get_local_linear_id()][1]; + }); + }, + exp_ext::property::node::depends_on{NodeA}); + + auto GraphExec = Graph.finalize(exp_ext::property::graph::updatable{}); + + // Submit the graph before the update and save the results. + Queue.submit([&](handler &CGH) { CGH.ext_oneapi_graph(GraphExec); }); + Queue.wait_and_throw(); + Queue.copy(PtrA, HostDataBeforeUpdate.data(), Size); + Queue.wait_and_throw(); + + DynLocalAccessor.update(range<2>{LocalMemSize * 2, 2}); + NodeA.update_nd_range(nd_range<1>(Size, LocalMemSize * 2)); + NodeB.update_nd_range(nd_range<1>(Size, LocalMemSize * 2)); + + GraphExec.update(NodeA); + GraphExec.update(NodeB); + + // Submit the graph after the update and save the results. + Queue.submit([&](handler &CGH) { CGH.ext_oneapi_graph(GraphExec); }); + Queue.wait_and_throw(); + Queue.copy(PtrA, HostDataAfterUpdate.data(), Size); + Queue.wait_and_throw(); + + for (size_t i = 0; i < Size; i++) { + int modI = i % LocalMemSize; + T Ref = (modI * 2) + (modI * 4); + assert(check_value(i, Ref, HostDataBeforeUpdate[i], "PtrA Before Update")); + } + + for (size_t i = 0; i < Size; i++) { + int modI = i % (LocalMemSize * 2); + T Ref = (modI * 2) + (modI * 4); + assert(check_value(i, Ref, HostDataAfterUpdate[i], "PtrA After Update")); + } + + free(PtrA, Queue); + + return 0; +} diff --git a/sycl/unittests/Extensions/CommandGraph/Update.cpp b/sycl/unittests/Extensions/CommandGraph/Update.cpp index b943b9c43dd98..c35659203b81c 100644 --- a/sycl/unittests/Extensions/CommandGraph/Update.cpp +++ b/sycl/unittests/Extensions/CommandGraph/Update.cpp @@ -40,6 +40,43 @@ TEST_F(CommandGraphTest, DynamicParamRegister) { }); } +TEST_F(CommandGraphTest, DynamicLocalAccessorRegister) { + // Check that registering a dynamic local accessor with a node from a graph + // that was not passed to its constructor throws. + experimental::dynamic_local_accessor DynamicLocalAcc(Graph, 10); + + auto OtherGraph = + experimental::command_graph(Queue.get_context(), Queue.get_device()); + auto Node = OtherGraph.add([&](sycl::handler &cgh) { + // This should throw since OtherGraph is not associated with DynamicParam + EXPECT_ANY_THROW(cgh.set_arg(0, DynamicLocalAcc)); + cgh.single_task>([]() {}); + }); +} + +TEST_F(CommandGraphTest, DynamicLocalAccessorNoGraph) { + // Check that using a dynamic local accessor in an eager sycl submission + // throws an exception. + experimental::dynamic_local_accessor DynamicLocalAcc(Graph, 10); + + Queue.submit([&](sycl::handler &cgh) { + EXPECT_ANY_THROW(cgh.set_arg(0, DynamicLocalAcc)); + cgh.single_task>([]() {}); + }); +} + +TEST_F(CommandGraphTest, DynamicLocalAccessorRecordingQueue) { + // Check that using a dynamic local accessor with a recording queue + // throws an exception. + experimental::dynamic_local_accessor DynamicLocalAcc(Graph, 10); + Graph.begin_recording(Queue); + + Queue.submit([&](sycl::handler &cgh) { + EXPECT_ANY_THROW(cgh.set_arg(0, DynamicLocalAcc)); + cgh.single_task>([]() {}); + }); +} + TEST_F(CommandGraphTest, UpdateNodeNotInGraph) { // Check that updating a graph with a node which is not part of that graph is // an error. From 06197104502d90a7dc8c10baed333f00a583c273 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?F=C3=A1bio=20Mestre?= Date: Tue, 14 Jan 2025 15:30:01 +0000 Subject: [PATCH 02/10] Fix UR cuda adapter to update indices correctly --- sycl/cmake/modules/FetchUnifiedRuntime.cmake | 2 +- sycl/cmake/modules/UnifiedRuntimeTag.cmake | 12 ++++++------ 2 files changed, 7 insertions(+), 7 deletions(-) diff --git a/sycl/cmake/modules/FetchUnifiedRuntime.cmake b/sycl/cmake/modules/FetchUnifiedRuntime.cmake index 72841724fa01d..202bfdcdc19e3 100644 --- a/sycl/cmake/modules/FetchUnifiedRuntime.cmake +++ b/sycl/cmake/modules/FetchUnifiedRuntime.cmake @@ -116,7 +116,7 @@ if(SYCL_UR_USE_FETCH_CONTENT) CACHE PATH "Path to external '${name}' adapter source dir" FORCE) endfunction() - set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git") + set(UNIFIED_RUNTIME_REPO "https://github.com/Bensuo/unified-runtime.git") include(${CMAKE_CURRENT_SOURCE_DIR}/cmake/modules/UnifiedRuntimeTag.cmake) set(UMF_BUILD_EXAMPLES OFF CACHE INTERNAL "EXAMPLES") diff --git a/sycl/cmake/modules/UnifiedRuntimeTag.cmake b/sycl/cmake/modules/UnifiedRuntimeTag.cmake index 526683c9cdf97..a50266c084229 100644 --- a/sycl/cmake/modules/UnifiedRuntimeTag.cmake +++ b/sycl/cmake/modules/UnifiedRuntimeTag.cmake @@ -1,7 +1,7 @@ -# commit 9e48f543b8dd39d45563169433bb529583625dfe -# Merge: 6a3fece6 1a1108b3 +# commit 6d4eec8cdcfe8a5d359ed05092797c429c2ca878 +# Merge: 40d28e7bd84a 800b452d67c5 # Author: Martin Grant -# Date: Wed Jan 15 14:33:29 2025 +0000 -# Merge pull request #2540 from martygrant/martin/program-info-unswitch -# Move urProgramGetInfo success test from a switch to individual tests. -set(UNIFIED_RUNTIME_TAG 9e48f543b8dd39d45563169433bb529583625dfe) +# Date: Thu Dec 12 16:00:13 2024 +0000 +# Merge pull request #2272 from martygrant/martin/virtual-memory-cts-spec-gap +# Improvements to align CTS and Spec for Virtual Memory +set(UNIFIED_RUNTIME_TAG bf6b6f9df5cd7c1e3dda4af8e4b3546c7109f24f) From 01ecad01d45a268c03444794c036774aec61b5b9 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?F=C3=A1bio=20Mestre?= Date: Tue, 14 Jan 2025 15:49:13 +0000 Subject: [PATCH 03/10] Fix formatting --- .../sycl/ext/oneapi/experimental/graph.hpp | 16 ++++++++++------ 1 file changed, 10 insertions(+), 6 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/graph.hpp b/sycl/include/sycl/ext/oneapi/experimental/graph.hpp index 898e0d06e595e..01ac53cb185a9 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/graph.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/graph.hpp @@ -17,7 +17,7 @@ #ifdef __INTEL_PREVIEW_BREAKING_CHANGES #include #endif -#include // for device +#include // for device #include // for graph properties classes #include // for range, nd_range #include // for is_property, is_property_of @@ -447,7 +447,6 @@ class command_graph namespace detail { class __SYCL_EXPORT dynamic_parameter_base { public: - dynamic_parameter_base( sycl::ext::oneapi::experimental::command_graph Graph); @@ -467,9 +466,11 @@ class __SYCL_EXPORT dynamic_parameter_base { void updateAccessor(const sycl::detail::AccessorBaseHost *Acc); - sycl::detail::LocalAccessorImplPtr getLocalAccessor(handler* Handler); + sycl::detail::LocalAccessorImplPtr getLocalAccessor(handler *Handler); - void registerLocalAccessor(sycl::detail::LocalAccessorBaseHost* LocalAccBaseHost, handler* Handler); + void + registerLocalAccessor(sycl::detail::LocalAccessorBaseHost *LocalAccBaseHost, + handler *Handler); void updateLocalAccessor(range<3> NewAllocationSize); @@ -530,14 +531,17 @@ class dynamic_local_accessor : public detail::dynamic_parameter_base { #ifndef __SYCL_DEVICE_ONLY__ ::sycl::detail::LocalAccessorImplPtr BaseLocalAcc = getLocalAccessor(&CGH); if (BaseLocalAcc) { - return sycl::detail::createSyclObjFromImpl>(BaseLocalAcc); + return sycl::detail::createSyclObjFromImpl< + local_accessor>(BaseLocalAcc); } else { local_accessor LocalAccessor(AllocationSize, CGH); registerLocalAccessor( - static_cast(&LocalAccessor), &CGH); + static_cast(&LocalAccessor), + &CGH); return LocalAccessor; } #else + (void)CGH; return local_accessor(); #endif }; From 097261f3d981ad65e55d4ab9d72e6f8428a95411 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?F=C3=A1bio=20Mestre?= Date: Tue, 14 Jan 2025 17:27:05 +0000 Subject: [PATCH 04/10] Fix abi and disable tests in L0 --- sycl/include/sycl/ext/oneapi/experimental/graph.hpp | 2 +- sycl/source/detail/graph_impl.hpp | 2 +- .../Graph/Update/dyn_cgf_dyn_local_accessor.cpp | 3 +++ .../Graph/Update/update_dynamic_local_accessor.cpp | 3 +++ ...pdate_dynamic_local_accessor_multiple_accessors.cpp | 3 +++ .../update_dynamic_local_accessor_multiple_nodes.cpp | 3 +++ sycl/test/abi/sycl_symbols_linux.dump | 10 ++++++---- 7 files changed, 20 insertions(+), 6 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/graph.hpp b/sycl/include/sycl/ext/oneapi/experimental/graph.hpp index 01ac53cb185a9..0437f59bc1ef0 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/graph.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/graph.hpp @@ -466,7 +466,7 @@ class __SYCL_EXPORT dynamic_parameter_base { void updateAccessor(const sycl::detail::AccessorBaseHost *Acc); - sycl::detail::LocalAccessorImplPtr getLocalAccessor(handler *Handler); + sycl::detail::LocalAccessorImplPtr getLocalAccessor(handler *Handler) const; void registerLocalAccessor(sycl::detail::LocalAccessorBaseHost *LocalAccBaseHost, diff --git a/sycl/source/detail/graph_impl.hpp b/sycl/source/detail/graph_impl.hpp index 993e1644a22fd..a2c06dd56330d 100644 --- a/sycl/source/detail/graph_impl.hpp +++ b/sycl/source/detail/graph_impl.hpp @@ -1492,7 +1492,7 @@ class dynamic_parameter_impl { /// @return returns the impl object for the local accessor that is associated /// with this handler. Or nullptr if no local accessor has been registered /// for this handler. - sycl::detail::LocalAccessorImplPtr getLocalAccessor(handler *Handler); + sycl::detail::LocalAccessorImplPtr getLocalAccessor(handler *Handler) const; /// Associates a local accessor with this dynamic local accessor for a /// specific handler. diff --git a/sycl/test-e2e/Graph/Update/dyn_cgf_dyn_local_accessor.cpp b/sycl/test-e2e/Graph/Update/dyn_cgf_dyn_local_accessor.cpp index 91cd54e7fb14c..8f70aff95c73e 100644 --- a/sycl/test-e2e/Graph/Update/dyn_cgf_dyn_local_accessor.cpp +++ b/sycl/test-e2e/Graph/Update/dyn_cgf_dyn_local_accessor.cpp @@ -5,6 +5,9 @@ // Extra run to check for immediate-command-list in Level Zero // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %} +// XFAIL: level_zero +// XFAIL-TRACKER: OFNAAO-422 + // Tests using dynamic command-group objects with dynamic local accessors. #include "../graph_common.hpp" diff --git a/sycl/test-e2e/Graph/Update/update_dynamic_local_accessor.cpp b/sycl/test-e2e/Graph/Update/update_dynamic_local_accessor.cpp index 2b485e53d9783..87a8b719360c9 100644 --- a/sycl/test-e2e/Graph/Update/update_dynamic_local_accessor.cpp +++ b/sycl/test-e2e/Graph/Update/update_dynamic_local_accessor.cpp @@ -5,6 +5,9 @@ // Extra run to check for immediate-command-list in Level Zero // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %} +// XFAIL: level_zero +// XFAIL-TRACKER: OFNAAO-422 + // Tests updating local accessor parameters. #include "../graph_common.hpp" diff --git a/sycl/test-e2e/Graph/Update/update_dynamic_local_accessor_multiple_accessors.cpp b/sycl/test-e2e/Graph/Update/update_dynamic_local_accessor_multiple_accessors.cpp index d3bf84c5f5f64..23aa083e16a10 100644 --- a/sycl/test-e2e/Graph/Update/update_dynamic_local_accessor_multiple_accessors.cpp +++ b/sycl/test-e2e/Graph/Update/update_dynamic_local_accessor_multiple_accessors.cpp @@ -5,6 +5,9 @@ // Extra run to check for immediate-command-list in Level Zero // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %} +// XFAIL: level_zero +// XFAIL-TRACKER: OFNAAO-422 + // Tests updating local accessor parameters using multiple dynamic local // accessors in the graph node. #include "../graph_common.hpp" diff --git a/sycl/test-e2e/Graph/Update/update_dynamic_local_accessor_multiple_nodes.cpp b/sycl/test-e2e/Graph/Update/update_dynamic_local_accessor_multiple_nodes.cpp index a10ae01207b2a..b0a5f5928ea35 100644 --- a/sycl/test-e2e/Graph/Update/update_dynamic_local_accessor_multiple_nodes.cpp +++ b/sycl/test-e2e/Graph/Update/update_dynamic_local_accessor_multiple_nodes.cpp @@ -5,6 +5,9 @@ // Extra run to check for immediate-command-list in Level Zero // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %} +// XFAIL: level_zero +// XFAIL-TRACKER: OFNAAO-422 + // Tests updating local 2D local accessor parameters in multiple graph nodes // simultaneously. Also tests using dynamic local accessors with // non-zero indices. diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 26a129e33ef85..53a36ffd08f31 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3072,7 +3072,12 @@ _ZN4sycl3_V13ext6oneapi12experimental6detail17build_from_sourceERNS0_13kernel_bu _ZN4sycl3_V13ext6oneapi12experimental6detail22dynamic_parameter_base11updateValueEPKNS3_14raw_kernel_argEm _ZN4sycl3_V13ext6oneapi12experimental6detail22dynamic_parameter_base11updateValueEPKvm _ZN4sycl3_V13ext6oneapi12experimental6detail22dynamic_parameter_base14updateAccessorEPKNS0_6detail16AccessorBaseHostE +_ZN4sycl3_V13ext6oneapi12experimental6detail22dynamic_parameter_base16getLocalAccessorEPNS0_7handlerE +_ZN4sycl3_V13ext6oneapi12experimental6detail22dynamic_parameter_base19updateLocalAccessorENS0_5rangeILi3EEE +_ZN4sycl3_V13ext6oneapi12experimental6detail22dynamic_parameter_base21registerLocalAccessorEPNS0_6detail21LocalAccessorBaseHostEPNS0_7handlerE +_ZN4sycl3_V13ext6oneapi12experimental6detail22dynamic_parameter_baseC1ENS3_13command_graphILNS3_11graph_stateE0EEE _ZN4sycl3_V13ext6oneapi12experimental6detail22dynamic_parameter_baseC1ENS3_13command_graphILNS3_11graph_stateE0EEEmPKv +_ZN4sycl3_V13ext6oneapi12experimental6detail22dynamic_parameter_baseC2ENS3_13command_graphILNS3_11graph_stateE0EEE _ZN4sycl3_V13ext6oneapi12experimental6detail22dynamic_parameter_baseC2ENS3_13command_graphILNS3_11graph_stateE0EEEmPKv _ZN4sycl3_V13ext6oneapi12experimental6detail24executable_command_graph12finalizeImplEv _ZN4sycl3_V13ext6oneapi12experimental6detail24executable_command_graph6updateERKNS3_13command_graphILNS3_11graph_stateE0EEE @@ -3086,7 +3091,6 @@ _ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graph13end_reco _ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graph15begin_recordingERKSt6vectorINS0_5queueESaIS7_EERKNS0_13property_listE _ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graph15begin_recordingERNS0_5queueERKNS0_13property_listE _ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graph24addGraphLeafDependenciesENS3_4nodeE -_ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graph27checkNodePropertiesAndThrowERKNS0_13property_listE _ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graph7addImplERKSt6vectorINS3_4nodeESaIS7_EE _ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graph7addImplERNS3_21dynamic_command_groupERKSt6vectorINS3_4nodeESaIS9_EE _ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graph7addImplESt8functionIFvRNS0_7handlerEEERKSt6vectorINS3_4nodeESaISC_EE @@ -3248,7 +3252,6 @@ _ZN4sycl3_V16detail13select_deviceERKSt8functionIFiRKNS0_6deviceEEE _ZN4sycl3_V16detail13select_deviceERKSt8functionIFiRKNS0_6deviceEEERKNS0_7contextE _ZN4sycl3_V16detail14SubmissionInfo14SecondaryQueueEv _ZN4sycl3_V16detail14SubmissionInfo17PostProcessorFuncEv -_ZN4sycl3_V16detail14SubmissionInfo9EventModeEv _ZN4sycl3_V16detail14SubmissionInfoC1Ev _ZN4sycl3_V16detail14SubmissionInfoC2Ev _ZN4sycl3_V16detail14addCounterInitERNS0_7handlerERSt10shared_ptrINS1_10queue_implEERS4_IiE @@ -3524,6 +3527,7 @@ _ZN4sycl3_V17handler20associateWithHandlerEPNS0_6detail30UnsampledImageAccessorB _ZN4sycl3_V17handler20memcpyToDeviceGlobalEPKvS3_bmm _ZN4sycl3_V17handler20setKernelCacheConfigENS1_23StableKernelCacheConfigE _ZN4sycl3_V17handler20setStateSpecConstSetEv +_ZN4sycl3_V17handler21setKernelWorkGroupMemEm _ZN4sycl3_V17handler21setUserFacingNodeTypeENS0_3ext6oneapi12experimental9node_typeE _ZN4sycl3_V17handler22ext_oneapi_fill2d_implEPvmPKvmmm _ZN4sycl3_V17handler22memcpyFromDeviceGlobalEPvPKvbmm @@ -3531,7 +3535,6 @@ _ZN4sycl3_V17handler22setHandlerKernelBundleENS0_6kernelE _ZN4sycl3_V17handler22setHandlerKernelBundleERKSt10shared_ptrINS0_6detail18kernel_bundle_implEE _ZN4sycl3_V17handler22setKernelClusterLaunchENS0_5rangeILi3EEEi _ZN4sycl3_V17handler22setKernelIsCooperativeEb -_ZN4sycl3_V17handler21setKernelWorkGroupMemEm _ZN4sycl3_V17handler24GetRangeRoundingSettingsERmS2_S2_ _ZN4sycl3_V17handler24ext_intel_read_host_pipeENS0_6detail11string_viewEPvmb _ZN4sycl3_V17handler24ext_oneapi_memcpy2d_implEPvmPKvmmm @@ -3725,7 +3728,6 @@ _ZNK4sycl3_V16detail12buffer_plain22get_allocator_internalEv _ZNK4sycl3_V16detail12buffer_plain7getSizeEv _ZNK4sycl3_V16detail14SubmissionInfo14SecondaryQueueEv _ZNK4sycl3_V16detail14SubmissionInfo17PostProcessorFuncEv -_ZNK4sycl3_V16detail14SubmissionInfo9EventModeEv _ZNK4sycl3_V16detail16AccessorBaseHost11getElemSizeEv _ZNK4sycl3_V16detail16AccessorBaseHost11getPropListEv _ZNK4sycl3_V16detail16AccessorBaseHost13isPlaceholderEv From 5b1a33ebccb18b60e5e3f9af90de4c2ee7ffb865 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?F=C3=A1bio=20Mestre?= Date: Tue, 14 Jan 2025 19:03:00 +0000 Subject: [PATCH 05/10] Address review comments --- .../sycl/ext/oneapi/experimental/graph.hpp | 10 +++++----- sycl/include/sycl/handler.hpp | 14 +++++++------- sycl/source/detail/graph_impl.cpp | 18 +++++++++++------- sycl/test/abi/sycl_symbols_linux.dump | 13 ++++++++----- 4 files changed, 31 insertions(+), 24 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/graph.hpp b/sycl/include/sycl/ext/oneapi/experimental/graph.hpp index 0437f59bc1ef0..f72ed979923bc 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/graph.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/graph.hpp @@ -449,12 +449,13 @@ class __SYCL_EXPORT dynamic_parameter_base { public: dynamic_parameter_base( sycl::ext::oneapi::experimental::command_graph - Graph); + Graph, + const property_list &PropList); dynamic_parameter_base( sycl::ext::oneapi::experimental::command_graph Graph, - size_t ParamSize, const void *Data); + size_t ParamSize, const void *Data, const property_list &PropList); protected: void updateValue(const void *NewValue, size_t Size); @@ -518,9 +519,8 @@ class dynamic_local_accessor : public detail::dynamic_parameter_base { dynamic_local_accessor(command_graph Graph, range AllocationSize, const property_list &PropList = {}) - : detail::dynamic_parameter_base(Graph), AllocationSize(AllocationSize) { - (void)PropList; - } + : detail::dynamic_parameter_base(Graph, PropList), + AllocationSize(AllocationSize) {} void update(range NewAllocationSize) { detail::dynamic_parameter_base::updateLocalAccessor( diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 4b6cc89c14ed9..4a02313c4c4a3 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -1850,21 +1850,21 @@ class __SYCL_EXPORT handler { // set_arg for graph dynamic_parameters template - void set_arg(int argIndex, - ext::oneapi::experimental::dynamic_parameter &dynamicParam) { - setArgHelper(argIndex, dynamicParam); + void set_arg(int ArgIndex, + ext::oneapi::experimental::dynamic_parameter &DynamicParam) { + setArgHelper(ArgIndex, DynamicParam); } template - void set_arg(int argIndex, + void set_arg(int ArgIndex, ext::oneapi::experimental::dynamic_local_accessor &DynamicLocalAccessor) { - setArgHelper(argIndex, DynamicLocalAccessor); + setArgHelper(ArgIndex, DynamicLocalAccessor); } // set_arg for the raw_kernel_arg extension type. - void set_arg(int argIndex, ext::oneapi::experimental::raw_kernel_arg &&Arg) { - setArgHelper(argIndex, std::move(Arg)); + void set_arg(int ArgIndex, ext::oneapi::experimental::raw_kernel_arg &&Arg) { + setArgHelper(ArgIndex, std::move(Arg)); } /// Sets arguments for OpenCL interoperability kernels. diff --git a/sycl/source/detail/graph_impl.cpp b/sycl/source/detail/graph_impl.cpp index 9baccd6b504fb..d242baefe0e0b 100644 --- a/sycl/source/detail/graph_impl.cpp +++ b/sycl/source/detail/graph_impl.cpp @@ -1894,15 +1894,19 @@ void executable_command_graph::update(const std::vector &Nodes) { } dynamic_parameter_base::dynamic_parameter_base( - command_graph Graph, size_t ParamSize, - const void *Data) + command_graph Graph, const property_list &PropList) : impl(std::make_shared( - sycl::detail::getSyclObjImpl(Graph), ParamSize, Data)) {} + sycl::detail::getSyclObjImpl(Graph))) { + checkGraphPropertiesAndThrow(PropList); +} dynamic_parameter_base::dynamic_parameter_base( - command_graph Graph) + command_graph Graph, size_t ParamSize, + const void *Data, const property_list &PropList) : impl(std::make_shared( - sycl::detail::getSyclObjImpl(Graph))) {} + sycl::detail::getSyclObjImpl(Graph), ParamSize, Data)) { + checkGraphPropertiesAndThrow(PropList); +} void dynamic_parameter_base::updateValue(const void *NewValue, size_t Size) { impl->updateValue(NewValue, Size); @@ -1919,7 +1923,7 @@ void dynamic_parameter_base::updateAccessor( } sycl::detail::LocalAccessorImplPtr -dynamic_parameter_base::getLocalAccessor(handler *Handler) { +dynamic_parameter_base::getLocalAccessor(handler *Handler) const { return impl->getLocalAccessor(Handler); } @@ -1988,7 +1992,7 @@ void dynamic_parameter_impl::updateAccessor( } sycl::detail::LocalAccessorImplPtr -dynamic_parameter_impl::getLocalAccessor(handler *Handler) { +dynamic_parameter_impl::getLocalAccessor(handler *Handler) const { auto HandlerImpl = sycl::detail::getSyclObjImpl(*Handler); auto FindLocalAcc = MHandlerToLocalAccMap.find(HandlerImpl); diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 53a36ffd08f31..49a3eec74b3af 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3072,13 +3072,12 @@ _ZN4sycl3_V13ext6oneapi12experimental6detail17build_from_sourceERNS0_13kernel_bu _ZN4sycl3_V13ext6oneapi12experimental6detail22dynamic_parameter_base11updateValueEPKNS3_14raw_kernel_argEm _ZN4sycl3_V13ext6oneapi12experimental6detail22dynamic_parameter_base11updateValueEPKvm _ZN4sycl3_V13ext6oneapi12experimental6detail22dynamic_parameter_base14updateAccessorEPKNS0_6detail16AccessorBaseHostE -_ZN4sycl3_V13ext6oneapi12experimental6detail22dynamic_parameter_base16getLocalAccessorEPNS0_7handlerE _ZN4sycl3_V13ext6oneapi12experimental6detail22dynamic_parameter_base19updateLocalAccessorENS0_5rangeILi3EEE _ZN4sycl3_V13ext6oneapi12experimental6detail22dynamic_parameter_base21registerLocalAccessorEPNS0_6detail21LocalAccessorBaseHostEPNS0_7handlerE -_ZN4sycl3_V13ext6oneapi12experimental6detail22dynamic_parameter_baseC1ENS3_13command_graphILNS3_11graph_stateE0EEE -_ZN4sycl3_V13ext6oneapi12experimental6detail22dynamic_parameter_baseC1ENS3_13command_graphILNS3_11graph_stateE0EEEmPKv -_ZN4sycl3_V13ext6oneapi12experimental6detail22dynamic_parameter_baseC2ENS3_13command_graphILNS3_11graph_stateE0EEE -_ZN4sycl3_V13ext6oneapi12experimental6detail22dynamic_parameter_baseC2ENS3_13command_graphILNS3_11graph_stateE0EEEmPKv +_ZN4sycl3_V13ext6oneapi12experimental6detail22dynamic_parameter_baseC1ENS3_13command_graphILNS3_11graph_stateE0EEERKNS0_13property_listE +_ZN4sycl3_V13ext6oneapi12experimental6detail22dynamic_parameter_baseC1ENS3_13command_graphILNS3_11graph_stateE0EEEmPKvRKNS0_13property_listE +_ZN4sycl3_V13ext6oneapi12experimental6detail22dynamic_parameter_baseC2ENS3_13command_graphILNS3_11graph_stateE0EEERKNS0_13property_listE +_ZN4sycl3_V13ext6oneapi12experimental6detail22dynamic_parameter_baseC2ENS3_13command_graphILNS3_11graph_stateE0EEEmPKvRKNS0_13property_listE _ZN4sycl3_V13ext6oneapi12experimental6detail24executable_command_graph12finalizeImplEv _ZN4sycl3_V13ext6oneapi12experimental6detail24executable_command_graph6updateERKNS3_13command_graphILNS3_11graph_stateE0EEE _ZN4sycl3_V13ext6oneapi12experimental6detail24executable_command_graph6updateERKNS3_4nodeE @@ -3091,6 +3090,7 @@ _ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graph13end_reco _ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graph15begin_recordingERKSt6vectorINS0_5queueESaIS7_EERKNS0_13property_listE _ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graph15begin_recordingERNS0_5queueERKNS0_13property_listE _ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graph24addGraphLeafDependenciesENS3_4nodeE +_ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graph27checkNodePropertiesAndThrowERKNS0_13property_listE _ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graph7addImplERKSt6vectorINS3_4nodeESaIS7_EE _ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graph7addImplERNS3_21dynamic_command_groupERKSt6vectorINS3_4nodeESaIS9_EE _ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graph7addImplESt8functionIFvRNS0_7handlerEEERKSt6vectorINS3_4nodeESaISC_EE @@ -3252,6 +3252,7 @@ _ZN4sycl3_V16detail13select_deviceERKSt8functionIFiRKNS0_6deviceEEE _ZN4sycl3_V16detail13select_deviceERKSt8functionIFiRKNS0_6deviceEEERKNS0_7contextE _ZN4sycl3_V16detail14SubmissionInfo14SecondaryQueueEv _ZN4sycl3_V16detail14SubmissionInfo17PostProcessorFuncEv +_ZN4sycl3_V16detail14SubmissionInfo9EventModeEv _ZN4sycl3_V16detail14SubmissionInfoC1Ev _ZN4sycl3_V16detail14SubmissionInfoC2Ev _ZN4sycl3_V16detail14addCounterInitERNS0_7handlerERSt10shared_ptrINS1_10queue_implEERS4_IiE @@ -3622,6 +3623,7 @@ _ZNK4sycl3_V13ext6oneapi12experimental21dynamic_command_group16get_active_indexE _ZNK4sycl3_V13ext6oneapi12experimental4node14get_successorsEv _ZNK4sycl3_V13ext6oneapi12experimental4node16get_predecessorsEv _ZNK4sycl3_V13ext6oneapi12experimental4node8get_typeEv +_ZNK4sycl3_V13ext6oneapi12experimental6detail22dynamic_parameter_base16getLocalAccessorEPNS0_7handlerE _ZNK4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graph11print_graphENS0_6detail11string_viewEb _ZNK4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graph11print_graphENSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEb _ZNK4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graph14get_root_nodesEv @@ -3728,6 +3730,7 @@ _ZNK4sycl3_V16detail12buffer_plain22get_allocator_internalEv _ZNK4sycl3_V16detail12buffer_plain7getSizeEv _ZNK4sycl3_V16detail14SubmissionInfo14SecondaryQueueEv _ZNK4sycl3_V16detail14SubmissionInfo17PostProcessorFuncEv +_ZNK4sycl3_V16detail14SubmissionInfo9EventModeEv _ZNK4sycl3_V16detail16AccessorBaseHost11getElemSizeEv _ZNK4sycl3_V16detail16AccessorBaseHost11getPropListEv _ZNK4sycl3_V16detail16AccessorBaseHost13isPlaceholderEv From c848b8019620b94c792a92fe1ca0caaf0612cb5b Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?F=C3=A1bio=20Mestre?= Date: Wed, 15 Jan 2025 12:47:21 +0000 Subject: [PATCH 06/10] Fix build issue --- sycl/include/sycl/ext/oneapi/experimental/graph.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/graph.hpp b/sycl/include/sycl/ext/oneapi/experimental/graph.hpp index f72ed979923bc..0f859ffdf8053 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/graph.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/graph.hpp @@ -499,7 +499,7 @@ class dynamic_parameter : public detail::dynamic_parameter_base { /// @param Param A reference value for this parameter used for CTAD. dynamic_parameter(experimental::command_graph Graph, const ValueT &Param) - : detail::dynamic_parameter_base(Graph, sizeof(ValueT), &Param) {} + : detail::dynamic_parameter_base(Graph, sizeof(ValueT), &Param, {}) {} /// Updates this dynamic parameter and all registered nodes with a new value. /// @param NewValue The new value for the parameter. From 80c211f21f62194b5494d033a244f087292ddf5c Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?F=C3=A1bio=20Mestre?= Date: Wed, 15 Jan 2025 15:29:44 +0000 Subject: [PATCH 07/10] Add windows symbols --- sycl/test/abi/sycl_symbols_windows.dump | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index a439081b1f382..d3d17c4192369 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -332,7 +332,11 @@ ??0device_selector@_V1@sycl@@QEAA@XZ ??0dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@QEAA@$$QEAV0123456@@Z ??0dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBV0123456@@Z -??0dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@QEAA@V?$command_graph@$0A@@23456@_KPEBX@Z +?updateLocalAccessor@dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@IEAAXV?$range@$02@67@@Z +??0dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@QEAA@V?$command_graph@$0A@@23456@AEBVproperty_list@56@@Z +?registerLocalAccessor@dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@IEAAXPEAVLocalAccessorBaseHost@267@PEAVhandler@67@@Z +??0dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@QEAA@V?$command_graph@$0A@@23456@_KPEBXAEBVproperty_list@56@@Z +?getLocalAccessor@dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@IEBA?AV?$shared_ptr@VLocalAccessorImplHost@detail@_V1@sycl@@@std@@PEAVhandler@67@@Z ?get_active_index@dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEBA_KXZ ??4dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEAAAEAV012345@$$QEAV012345@@Z ?addImpl@modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@IEAA?AVnode@34567@AEAVdynamic_command_group@34567@AEBV?$vector@Vnode@experimental@oneapi@ext@_V1@sycl@@V?$allocator@Vnode@experimental@oneapi@ext@_V1@sycl@@@std@@@std@@@Z From 5c5899533c62c2422ab7d981a1a5d312b3f8d5f5 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?F=C3=A1bio=20Mestre?= Date: Wed, 15 Jan 2025 17:40:52 +0000 Subject: [PATCH 08/10] Make dynamic_parameter_base Proplist parameter have a default --- sycl/include/sycl/ext/oneapi/experimental/graph.hpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/graph.hpp b/sycl/include/sycl/ext/oneapi/experimental/graph.hpp index 0f859ffdf8053..f4e320560a65d 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/graph.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/graph.hpp @@ -450,12 +450,12 @@ class __SYCL_EXPORT dynamic_parameter_base { dynamic_parameter_base( sycl::ext::oneapi::experimental::command_graph Graph, - const property_list &PropList); + const property_list &PropList = {}); dynamic_parameter_base( sycl::ext::oneapi::experimental::command_graph Graph, - size_t ParamSize, const void *Data, const property_list &PropList); + size_t ParamSize, const void *Data, const property_list &PropList = {}); protected: void updateValue(const void *NewValue, size_t Size); @@ -499,7 +499,7 @@ class dynamic_parameter : public detail::dynamic_parameter_base { /// @param Param A reference value for this parameter used for CTAD. dynamic_parameter(experimental::command_graph Graph, const ValueT &Param) - : detail::dynamic_parameter_base(Graph, sizeof(ValueT), &Param, {}) {} + : detail::dynamic_parameter_base(Graph, sizeof(ValueT), &Param) {} /// Updates this dynamic parameter and all registered nodes with a new value. /// @param NewValue The new value for the parameter. From 2c037ec8d5144a5a075f31de003502bf71a7735f Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?F=C3=A1bio=20Mestre?= Date: Wed, 15 Jan 2025 18:03:26 +0000 Subject: [PATCH 09/10] Remove PropList from existing constructor --- sycl/include/sycl/ext/oneapi/experimental/graph.hpp | 4 ++-- sycl/source/detail/graph_impl.cpp | 6 ++---- 2 files changed, 4 insertions(+), 6 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/graph.hpp b/sycl/include/sycl/ext/oneapi/experimental/graph.hpp index f4e320560a65d..6bdd9444875e5 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/graph.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/graph.hpp @@ -450,12 +450,12 @@ class __SYCL_EXPORT dynamic_parameter_base { dynamic_parameter_base( sycl::ext::oneapi::experimental::command_graph Graph, - const property_list &PropList = {}); + const property_list &PropList); dynamic_parameter_base( sycl::ext::oneapi::experimental::command_graph Graph, - size_t ParamSize, const void *Data, const property_list &PropList = {}); + size_t ParamSize, const void *Data); protected: void updateValue(const void *NewValue, size_t Size); diff --git a/sycl/source/detail/graph_impl.cpp b/sycl/source/detail/graph_impl.cpp index d242baefe0e0b..7262fa2088f1e 100644 --- a/sycl/source/detail/graph_impl.cpp +++ b/sycl/source/detail/graph_impl.cpp @@ -1902,11 +1902,9 @@ dynamic_parameter_base::dynamic_parameter_base( dynamic_parameter_base::dynamic_parameter_base( command_graph Graph, size_t ParamSize, - const void *Data, const property_list &PropList) + const void *Data) : impl(std::make_shared( - sycl::detail::getSyclObjImpl(Graph), ParamSize, Data)) { - checkGraphPropertiesAndThrow(PropList); -} + sycl::detail::getSyclObjImpl(Graph), ParamSize, Data)) {} void dynamic_parameter_base::updateValue(const void *NewValue, size_t Size) { impl->updateValue(NewValue, Size); From c77c115da360d3d3a86f93add9bf0d36144d3613 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?F=C3=A1bio=20Mestre?= Date: Wed, 15 Jan 2025 18:50:38 +0000 Subject: [PATCH 10/10] Update symbols --- sycl/test/abi/sycl_symbols_linux.dump | 4 ++-- sycl/test/abi/sycl_symbols_windows.dump | 2 +- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 49a3eec74b3af..cf6c9f460d0cf 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3075,9 +3075,9 @@ _ZN4sycl3_V13ext6oneapi12experimental6detail22dynamic_parameter_base14updateAcce _ZN4sycl3_V13ext6oneapi12experimental6detail22dynamic_parameter_base19updateLocalAccessorENS0_5rangeILi3EEE _ZN4sycl3_V13ext6oneapi12experimental6detail22dynamic_parameter_base21registerLocalAccessorEPNS0_6detail21LocalAccessorBaseHostEPNS0_7handlerE _ZN4sycl3_V13ext6oneapi12experimental6detail22dynamic_parameter_baseC1ENS3_13command_graphILNS3_11graph_stateE0EEERKNS0_13property_listE -_ZN4sycl3_V13ext6oneapi12experimental6detail22dynamic_parameter_baseC1ENS3_13command_graphILNS3_11graph_stateE0EEEmPKvRKNS0_13property_listE +_ZN4sycl3_V13ext6oneapi12experimental6detail22dynamic_parameter_baseC1ENS3_13command_graphILNS3_11graph_stateE0EEEmPKv _ZN4sycl3_V13ext6oneapi12experimental6detail22dynamic_parameter_baseC2ENS3_13command_graphILNS3_11graph_stateE0EEERKNS0_13property_listE -_ZN4sycl3_V13ext6oneapi12experimental6detail22dynamic_parameter_baseC2ENS3_13command_graphILNS3_11graph_stateE0EEEmPKvRKNS0_13property_listE +_ZN4sycl3_V13ext6oneapi12experimental6detail22dynamic_parameter_baseC2ENS3_13command_graphILNS3_11graph_stateE0EEEmPKv _ZN4sycl3_V13ext6oneapi12experimental6detail24executable_command_graph12finalizeImplEv _ZN4sycl3_V13ext6oneapi12experimental6detail24executable_command_graph6updateERKNS3_13command_graphILNS3_11graph_stateE0EEE _ZN4sycl3_V13ext6oneapi12experimental6detail24executable_command_graph6updateERKNS3_4nodeE diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index d3d17c4192369..e7fbd2ba781ae 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -335,7 +335,7 @@ ?updateLocalAccessor@dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@IEAAXV?$range@$02@67@@Z ??0dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@QEAA@V?$command_graph@$0A@@23456@AEBVproperty_list@56@@Z ?registerLocalAccessor@dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@IEAAXPEAVLocalAccessorBaseHost@267@PEAVhandler@67@@Z -??0dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@QEAA@V?$command_graph@$0A@@23456@_KPEBXAEBVproperty_list@56@@Z +??0dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@QEAA@V?$command_graph@$0A@@23456@_KPEBX@Z ?getLocalAccessor@dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@IEBA?AV?$shared_ptr@VLocalAccessorImplHost@detail@_V1@sycl@@@std@@PEAVhandler@67@@Z ?get_active_index@dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEBA_KXZ ??4dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEAAAEAV012345@$$QEAV012345@@Z