diff --git a/sycl/include/CL/sycl/detail/cg.hpp b/sycl/include/CL/sycl/detail/cg.hpp index 0b2ebddbf3a24..ba70ddefcee2e 100644 --- a/sycl/include/CL/sycl/detail/cg.hpp +++ b/sycl/include/CL/sycl/detail/cg.hpp @@ -248,6 +248,7 @@ class CGExecKernel : public CG { std::string MKernelName; detail::OSModuleHandle MOSModuleHandle; std::vector> MStreams; + std::vector> MAuxiliaryResources; CGExecKernel(NDRDescT NDRDesc, std::unique_ptr HKernel, std::shared_ptr SyclKernel, @@ -259,6 +260,7 @@ class CGExecKernel : public CG { std::vector Args, std::string KernelName, detail::OSModuleHandle OSModuleHandle, std::vector> Streams, + std::vector> AuxiliaryResources, CGTYPE Type, detail::code_location loc = {}) : CG(Type, std::move(ArgsStorage), std::move(AccStorage), std::move(SharedPtrStorage), std::move(Requirements), @@ -266,7 +268,8 @@ class CGExecKernel : public CG { MNDRDesc(std::move(NDRDesc)), MHostKernel(std::move(HKernel)), MSyclKernel(std::move(SyclKernel)), MArgs(std::move(Args)), MKernelName(std::move(KernelName)), MOSModuleHandle(OSModuleHandle), - MStreams(std::move(Streams)) { + MStreams(std::move(Streams)), + MAuxiliaryResources(std::move(AuxiliaryResources)) { assert((getType() == RunOnHostIntel || getType() == Kernel) && "Wrong type of exec kernel CG."); } @@ -277,6 +280,10 @@ class CGExecKernel : public CG { return MStreams; } + std::vector> getAuxiliaryResources() const { + return MAuxiliaryResources; + } + std::shared_ptr getKernelBundle() { const std::shared_ptr> &ExtendedMembers = getExtendedMembers(); @@ -291,6 +298,9 @@ class CGExecKernel : public CG { void clearStreams() { MStreams.clear(); } bool hasStreams() { return !MStreams.empty(); } + + void clearAuxiliaryResources() { MAuxiliaryResources.clear(); } + bool hasAuxiliaryResources() { return !MAuxiliaryResources.empty(); } }; /// "Copy memory" command group class. diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp index aa0c80d546b72..3004f438e32b9 100644 --- a/sycl/include/CL/sycl/handler.hpp +++ b/sycl/include/CL/sycl/handler.hpp @@ -472,12 +472,9 @@ class __SYCL_EXPORT handler { /// Saves buffers created by handling reduction feature in handler. /// They are then forwarded to command group and destroyed only after /// the command group finishes the work on device/host. - /// The 'MSharedPtrStorage' suits that need. /// /// @param ReduObj is a pointer to object that must be stored. - void addReduction(const std::shared_ptr &ReduObj) { - MSharedPtrStorage.push_back(ReduObj); - } + void addReduction(const std::shared_ptr &ReduObj); ~handler() = default; @@ -1280,6 +1277,7 @@ class __SYCL_EXPORT handler { } std::shared_ptr getHandlerImpl() const; + std::shared_ptr evictHandlerImpl() const; void setStateExplicitKernelBundle(); void setStateSpecConstSet(); diff --git a/sycl/include/sycl/ext/oneapi/reduction.hpp b/sycl/include/sycl/ext/oneapi/reduction.hpp index 2944c8832275b..1b052f26108cb 100644 --- a/sycl/include/sycl/ext/oneapi/reduction.hpp +++ b/sycl/include/sycl/ext/oneapi/reduction.hpp @@ -718,6 +718,7 @@ class reduction_impl : private reduction_impl_base { auto RWReduVal = std::make_shared(MIdentity); CGH.addReduction(RWReduVal); MOutBufPtr = std::make_shared>(RWReduVal.get(), range<1>(1)); + MOutBufPtr->set_final_data(); CGH.addReduction(MOutBufPtr); return createHandlerWiredReadWriteAccessor(CGH, *MOutBufPtr); } @@ -728,6 +729,7 @@ class reduction_impl : private reduction_impl_base { auto CounterMem = std::make_shared(0); CGH.addReduction(CounterMem); auto CounterBuf = std::make_shared>(CounterMem.get(), 1); + CounterBuf->set_final_data(); CGH.addReduction(CounterBuf); return {*CounterBuf, CGH}; } diff --git a/sycl/source/detail/handler_impl.hpp b/sycl/source/detail/handler_impl.hpp index d4171e8d4d1d6..673650874bf7d 100644 --- a/sycl/source/detail/handler_impl.hpp +++ b/sycl/source/detail/handler_impl.hpp @@ -65,6 +65,9 @@ class handler_impl { /// equal to the queue associated with the handler if the corresponding /// submission is a fallback from a previous submission. std::shared_ptr MSubmissionSecondaryQueue; + + // Stores auxiliary resources used by internal operations. + std::vector> MAuxiliaryResources; }; } // namespace detail diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp index 0af040568d0a9..f9b36419f6294 100644 --- a/sycl/source/detail/scheduler/commands.cpp +++ b/sycl/source/detail/scheduler/commands.cpp @@ -1378,11 +1378,23 @@ std::vector ExecCGCommand::getStreams() const { return {}; } +std::vector> +ExecCGCommand::getAuxiliaryResources() const { + if (MCommandGroup->getType() == CG::Kernel) + return ((CGExecKernel *)MCommandGroup.get())->getAuxiliaryResources(); + return {}; +} + void ExecCGCommand::clearStreams() { if (MCommandGroup->getType() == CG::Kernel) ((CGExecKernel *)MCommandGroup.get())->clearStreams(); } +void ExecCGCommand::clearAuxiliaryResources() { + if (MCommandGroup->getType() == CG::Kernel) + ((CGExecKernel *)MCommandGroup.get())->clearAuxiliaryResources(); +} + cl_int UpdateHostRequirementCommand::enqueueImp() { waitForPreparedHostEvents(); std::vector EventImpls = MPreparedDepsEvents; @@ -1673,7 +1685,9 @@ ExecCGCommand::ExecCGCommand(std::unique_ptr CommandGroup, static_cast(MCommandGroup.get())->MQueue; MEvent->setNeedsCleanupAfterWait(true); } else if (MCommandGroup->getType() == CG::CGTYPE::Kernel && - (static_cast(MCommandGroup.get()))->hasStreams()) + (static_cast(MCommandGroup.get())->hasStreams() || + static_cast(MCommandGroup.get()) + ->hasAuxiliaryResources())) MEvent->setNeedsCleanupAfterWait(true); emitInstrumentationDataProxy(); @@ -2481,7 +2495,9 @@ bool ExecCGCommand::supportsPostEnqueueCleanup() const { return Command::supportsPostEnqueueCleanup() && (MCommandGroup->getType() != CG::CGTYPE::CodeplayHostTask) && (MCommandGroup->getType() != CG::CGTYPE::Kernel || - !(static_cast(MCommandGroup.get()))->hasStreams()); + (!static_cast(MCommandGroup.get())->hasStreams() && + !static_cast(MCommandGroup.get()) + ->hasAuxiliaryResources())); } } // namespace detail diff --git a/sycl/source/detail/scheduler/commands.hpp b/sycl/source/detail/scheduler/commands.hpp index 6684dd588bf14..0a4a55579c078 100644 --- a/sycl/source/detail/scheduler/commands.hpp +++ b/sycl/source/detail/scheduler/commands.hpp @@ -541,8 +541,10 @@ class ExecCGCommand : public Command { ExecCGCommand(std::unique_ptr CommandGroup, QueueImplPtr Queue); std::vector getStreams() const; + std::vector> getAuxiliaryResources() const; void clearStreams(); + void clearAuxiliaryResources(); void printDot(std::ostream &Stream) const final; void emitInstrumentationData() final; diff --git a/sycl/source/detail/scheduler/graph_builder.cpp b/sycl/source/detail/scheduler/graph_builder.cpp index ed2ee3e6f78dc..4899e079cbbc6 100644 --- a/sycl/source/detail/scheduler/graph_builder.cpp +++ b/sycl/source/detail/scheduler/graph_builder.cpp @@ -1045,7 +1045,8 @@ void Scheduler::GraphBuilder::decrementLeafCountersForRecord( void Scheduler::GraphBuilder::cleanupCommandsForRecord( MemObjRecord *Record, - std::vector> &StreamsToDeallocate) { + std::vector> &StreamsToDeallocate, + std::vector> &AuxResourcesToDeallocate) { std::vector &AllocaCommands = Record->MAllocaCommands; if (AllocaCommands.empty()) return; @@ -1097,10 +1098,19 @@ void Scheduler::GraphBuilder::cleanupCommandsForRecord( // Collect stream objects for a visited command. if (Cmd->getType() == Command::CommandType::RUN_CG) { auto ExecCmd = static_cast(Cmd); + + // Transfer ownership of stream implementations. std::vector> Streams = ExecCmd->getStreams(); ExecCmd->clearStreams(); StreamsToDeallocate.insert(StreamsToDeallocate.end(), Streams.begin(), Streams.end()); + + // Transfer ownership of auxiliary resources. + std::vector> AuxResources = + ExecCmd->getAuxiliaryResources(); + ExecCmd->clearAuxiliaryResources(); + AuxResourcesToDeallocate.insert(AuxResourcesToDeallocate.end(), + AuxResources.begin(), AuxResources.end()); } for (Command *UserCmd : Cmd->MUsers) @@ -1160,6 +1170,7 @@ void Scheduler::GraphBuilder::cleanupCommand(Command *Cmd) { if (ExecCGCmd->getCG().getType() == CG::CGTYPE::Kernel) { auto *ExecKernelCG = static_cast(&ExecCGCmd->getCG()); assert(!ExecKernelCG->hasStreams()); + assert(!ExecKernelCG->hasAuxiliaryResources()); } } #endif @@ -1191,7 +1202,8 @@ void Scheduler::GraphBuilder::cleanupCommand(Command *Cmd) { void Scheduler::GraphBuilder::cleanupFinishedCommands( Command *FinishedCmd, - std::vector> &StreamsToDeallocate) { + std::vector> &StreamsToDeallocate, + std::vector> &AuxResourcesToDeallocate) { assert(MCmdsToVisit.empty()); MCmdsToVisit.push(FinishedCmd); MVisitedCmds.clear(); @@ -1207,10 +1219,19 @@ void Scheduler::GraphBuilder::cleanupFinishedCommands( // Collect stream objects for a visited command. if (Cmd->getType() == Command::CommandType::RUN_CG) { auto ExecCmd = static_cast(Cmd); + + // Transfer ownership of stream implementations. std::vector> Streams = ExecCmd->getStreams(); ExecCmd->clearStreams(); StreamsToDeallocate.insert(StreamsToDeallocate.end(), Streams.begin(), Streams.end()); + + // Transfer ownership of auxiliary resources. + std::vector> AuxResources = + ExecCmd->getAuxiliaryResources(); + ExecCmd->clearAuxiliaryResources(); + AuxResourcesToDeallocate.insert(AuxResourcesToDeallocate.end(), + AuxResources.begin(), AuxResources.end()); } for (const DepDesc &Dep : Cmd->MDeps) { diff --git a/sycl/source/detail/scheduler/scheduler.cpp b/sycl/source/detail/scheduler/scheduler.cpp index 2eda15bdeba49..8354e9c8d7b0b 100644 --- a/sycl/source/detail/scheduler/scheduler.cpp +++ b/sycl/source/detail/scheduler/scheduler.cpp @@ -239,6 +239,11 @@ void Scheduler::cleanupFinishedCommands(EventImplPtr FinishedEvent) { // objects, this is needed to guarantee that streamed data is printed and // resources are released. std::vector> StreamsToDeallocate; + // Similar to streams, we also collect the auxiliary resources used by the + // commands. Cleanup will make sure the commands do not own the resources + // anymore, so we just need them to survive the graph lock then they can die + // as they go out of scope. + std::vector> AuxResourcesToDeallocate; { // Avoiding deadlock situation, where one thread is in the process of // enqueueing (with a locked mutex) a currently blocked task that waits for @@ -249,7 +254,8 @@ void Scheduler::cleanupFinishedCommands(EventImplPtr FinishedEvent) { // The command might have been cleaned up (and set to nullptr) by another // thread if (FinishedCmd) - MGraphBuilder.cleanupFinishedCommands(FinishedCmd, StreamsToDeallocate); + MGraphBuilder.cleanupFinishedCommands(FinishedCmd, StreamsToDeallocate, + AuxResourcesToDeallocate); } } deallocateStreams(StreamsToDeallocate); @@ -261,6 +267,11 @@ void Scheduler::removeMemoryObject(detail::SYCLMemObjI *MemObj) { // objects, this is needed to guarantee that streamed data is printed and // resources are released. std::vector> StreamsToDeallocate; + // Similar to streams, we also collect the auxiliary resources used by the + // commands. Cleanup will make sure the commands do not own the resources + // anymore, so we just need them to survive the graph lock then they can die + // as they go out of scope. + std::vector> AuxResourcesToDeallocate; { MemObjRecord *Record = nullptr; @@ -282,7 +293,8 @@ void Scheduler::removeMemoryObject(detail::SYCLMemObjI *MemObj) { WriteLockT Lock(MGraphLock, std::defer_lock); acquireWriteLock(Lock); MGraphBuilder.decrementLeafCountersForRecord(Record); - MGraphBuilder.cleanupCommandsForRecord(Record, StreamsToDeallocate); + MGraphBuilder.cleanupCommandsForRecord(Record, StreamsToDeallocate, + AuxResourcesToDeallocate); MGraphBuilder.removeRecordForMemObj(MemObj); } } diff --git a/sycl/source/detail/scheduler/scheduler.hpp b/sycl/source/detail/scheduler/scheduler.hpp index 18ed2f5004c06..39075dbcd2703 100644 --- a/sycl/source/detail/scheduler/scheduler.hpp +++ b/sycl/source/detail/scheduler/scheduler.hpp @@ -514,7 +514,8 @@ class Scheduler { /// (assuming that all its commands have been waited for). void cleanupFinishedCommands( Command *FinishedCmd, - std::vector> &); + std::vector> &, + std::vector> &); /// Reschedules the command passed using Queue provided. /// @@ -540,7 +541,8 @@ class Scheduler { /// Removes commands that use the given MemObjRecord from the graph. void cleanupCommandsForRecord( MemObjRecord *Record, - std::vector> &); + std::vector> &, + std::vector> &); /// Removes the MemObjRecord for the memory object passed. void removeRecordForMemObj(SYCLMemObjI *MemObject); diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index 5f702bcdde6ff..aa349c90a6f33 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -49,24 +49,40 @@ handler::handler(std::shared_ptr Queue, MSharedPtrStorage.push_back(std::move(ExtendedMembers)); } +static detail::ExtendedMemberT &getHandlerImplMember( + std::vector> &SharedPtrStorage) { + assert(!SharedPtrStorage.empty()); + std::shared_ptr> ExtendedMembersVec = + detail::convertToExtendedMembers(SharedPtrStorage[0]); + assert(ExtendedMembersVec->size() > 0); + auto &HandlerImplMember = (*ExtendedMembersVec)[0]; + assert(detail::ExtendedMembersType::HANDLER_IMPL == HandlerImplMember.MType); + return HandlerImplMember; +} + /// Gets the handler_impl at the start of the extended members. std::shared_ptr handler::getHandlerImpl() const { std::lock_guard Lock( detail::GlobalHandler::instance().getHandlerExtendedMembersMutex()); + return std::static_pointer_cast( + getHandlerImplMember(MSharedPtrStorage).MData); +} - assert(!MSharedPtrStorage.empty()); - - std::shared_ptr> ExtendedMembersVec = - detail::convertToExtendedMembers(MSharedPtrStorage[0]); - - assert(ExtendedMembersVec->size() > 0); - - auto HandlerImplMember = (*ExtendedMembersVec)[0]; +/// Gets the handler_impl at the start of the extended members and removes it. +std::shared_ptr handler::evictHandlerImpl() const { + std::lock_guard Lock( + detail::GlobalHandler::instance().getHandlerExtendedMembersMutex()); + auto &HandlerImplMember = getHandlerImplMember(MSharedPtrStorage); + auto Impl = + std::static_pointer_cast(HandlerImplMember.MData); - assert(detail::ExtendedMembersType::HANDLER_IMPL == HandlerImplMember.MType); + // Reset the data of the member. + // NOTE: We let it stay because removing the front can be expensive. This will + // be improved when the impl is made a member of handler. In fact eviction is + // likely to not be needed when that happens. + HandlerImplMember.MData.reset(); - return std::static_pointer_cast( - HandlerImplMember.MData); + return Impl; } // Sets the submission state to indicate that an explicit kernel bundle has been @@ -281,6 +297,10 @@ event handler::finalize() { return MLastEvent; } + // Evict handler_impl from extended members to make sure the command group + // does not keep it alive. + std::shared_ptr Impl = evictHandlerImpl(); + std::unique_ptr CommandGroup; switch (type) { case detail::CG::Kernel: @@ -293,7 +313,8 @@ event handler::finalize() { std::move(MArgsStorage), std::move(MAccStorage), std::move(MSharedPtrStorage), std::move(MRequirements), std::move(MEvents), std::move(MArgs), MKernelName, MOSModuleHandle, - std::move(MStreamStorage), MCGType, MCodeLoc)); + std::move(MStreamStorage), std::move(Impl->MAuxiliaryResources), + MCGType, MCodeLoc)); break; } case detail::CG::CodeplayInteropTask: @@ -382,6 +403,10 @@ event handler::finalize() { return MLastEvent; } +void handler::addReduction(const std::shared_ptr &ReduObj) { + getHandlerImpl()->MAuxiliaryResources.push_back(ReduObj); +} + void handler::associateWithHandler(detail::AccessorBaseHost *AccBase, access::target AccTarget) { detail::AccessorImplPtr AccImpl = detail::getSyclObjImpl(*AccBase); diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 79d58722b4f4f..5587ff14009b6 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3994,6 +3994,7 @@ _ZN2cl4sycl7handler10depends_onERKSt6vectorINS0_5eventESaIS3_EE _ZN2cl4sycl7handler10mem_adviseEPKvmi _ZN2cl4sycl7handler10processArgEPvRKNS0_6detail19kernel_param_kind_tEimRmb _ZN2cl4sycl7handler10processArgEPvRKNS0_6detail19kernel_param_kind_tEimRmbb +_ZN2cl4sycl7handler12addReductionERKSt10shared_ptrIKvE _ZN2cl4sycl7handler13getKernelNameB5cxx11Ev _ZN2cl4sycl7handler17use_kernel_bundleERKNS0_13kernel_bundleILNS0_12bundle_stateE2EEE _ZN2cl4sycl7handler18RangeRoundingTraceEv @@ -4390,6 +4391,7 @@ _ZNK2cl4sycl7context8get_infoILNS0_4info7contextE65552EEENS3_12param_traitsIS4_X _ZNK2cl4sycl7context8get_infoILNS0_4info7contextE65553EEENS3_12param_traitsIS4_XT_EE11return_typeEv _ZNK2cl4sycl7context9getNativeEv _ZNK2cl4sycl7handler14getHandlerImplEv +_ZNK2cl4sycl7handler16evictHandlerImplEv _ZNK2cl4sycl7handler27isStateExplicitKernelBundleEv _ZNK2cl4sycl7handler30getOrInsertHandlerKernelBundleEb _ZNK2cl4sycl7program10get_kernelENSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index 4c6d0e99c0c00..03196958fed0e 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -1754,6 +1754,7 @@ ?erfc@__host_std@cl@@YA?AVhalf@half_impl@detail@sycl@2@V34562@@Z ?erfc@__host_std@cl@@YAMM@Z ?erfc@__host_std@cl@@YANN@Z +?evictHandlerImpl@handler@sycl@cl@@AEBA?AV?$shared_ptr@Vhandler_impl@detail@sycl@cl@@@std@@XZ ?exp10@__host_std@cl@@YA?AV?$vec@M$00@sycl@2@V342@@Z ?exp10@__host_std@cl@@YA?AV?$vec@M$01@sycl@2@V342@@Z ?exp10@__host_std@cl@@YA?AV?$vec@M$02@sycl@2@V342@@Z diff --git a/sycl/unittests/program_manager/EliminatedArgMask.cpp b/sycl/unittests/program_manager/EliminatedArgMask.cpp index 5301ea986ad94..2fcb0750e13af 100644 --- a/sycl/unittests/program_manager/EliminatedArgMask.cpp +++ b/sycl/unittests/program_manager/EliminatedArgMask.cpp @@ -7,6 +7,7 @@ //===----------------------------------------------------------------------===// #include +#include #include #include #include @@ -126,6 +127,7 @@ class MockHandler : public sycl::handler { std::unique_ptr finalize() { auto CGH = static_cast(this); + std::shared_ptr Impl = evictHandlerImpl(); std::unique_ptr CommandGroup; switch (getType()) { case sycl::detail::CG::Kernel: { @@ -136,7 +138,7 @@ class MockHandler : public sycl::handler { std::move(CGH->MRequirements), std::move(CGH->MEvents), std::move(CGH->MArgs), std::move(CGH->MKernelName), std::move(CGH->MOSModuleHandle), std::move(CGH->MStreamStorage), - CGH->MCGType, CGH->MCodeLoc)); + std::move(Impl->MAuxiliaryResources), CGH->MCGType, CGH->MCodeLoc)); break; } default: diff --git a/sycl/unittests/scheduler/Regression.cpp b/sycl/unittests/scheduler/Regression.cpp index f0af4be0665eb..48f7f9c44bd21 100644 --- a/sycl/unittests/scheduler/Regression.cpp +++ b/sycl/unittests/scheduler/Regression.cpp @@ -86,6 +86,7 @@ TEST_F(SchedulerTest, CheckArgsBlobInPiEnqueueNativeKernelIsValid) { /*KernelName*/ "", /*OSModuleHandle*/ detail::OSUtil::ExeModuleHandle, /*Streams*/ {}, + /*AuxiliaryResources*/ {}, /*Type*/ detail::CG::RunOnHostIntel)}; context Ctx{Plt}; diff --git a/sycl/unittests/scheduler/SchedulerTestUtils.hpp b/sycl/unittests/scheduler/SchedulerTestUtils.hpp index 2f6073d1d3672..f4dcc4e0f215f 100644 --- a/sycl/unittests/scheduler/SchedulerTestUtils.hpp +++ b/sycl/unittests/scheduler/SchedulerTestUtils.hpp @@ -124,7 +124,9 @@ class MockScheduler : public cl::sycl::detail::Scheduler { void cleanupCommandsForRecord(cl::sycl::detail::MemObjRecord *Rec) { std::vector> StreamsToDeallocate; - MGraphBuilder.cleanupCommandsForRecord(Rec, StreamsToDeallocate); + std::vector> AuxiliaryResourcesToDeallocate; + MGraphBuilder.cleanupCommandsForRecord(Rec, StreamsToDeallocate, + AuxiliaryResourcesToDeallocate); } void addNodeToLeaves(cl::sycl::detail::MemObjRecord *Rec, diff --git a/sycl/unittests/scheduler/StreamInitDependencyOnHost.cpp b/sycl/unittests/scheduler/StreamInitDependencyOnHost.cpp index 478465603199a..ef8c4c0895df8 100644 --- a/sycl/unittests/scheduler/StreamInitDependencyOnHost.cpp +++ b/sycl/unittests/scheduler/StreamInitDependencyOnHost.cpp @@ -10,6 +10,7 @@ #include "SchedulerTestUtils.hpp" #include +#include #include #include @@ -44,6 +45,7 @@ class MockHandler : public sycl::handler { std::unique_ptr finalize() { auto CGH = static_cast(this); + std::shared_ptr Impl = evictHandlerImpl(); std::unique_ptr CommandGroup; switch (CGH->MCGType) { case detail::CG::Kernel: @@ -55,7 +57,7 @@ class MockHandler : public sycl::handler { std::move(CGH->MRequirements), std::move(CGH->MEvents), std::move(CGH->MArgs), std::move(CGH->MKernelName), std::move(CGH->MOSModuleHandle), std::move(CGH->MStreamStorage), - CGH->MCGType, CGH->MCodeLoc)); + std::move(Impl->MAuxiliaryResources), CGH->MCGType, CGH->MCodeLoc)); break; } default: