From 20ea2920c6ba345d804c7a9952f4eff1872d03fd Mon Sep 17 00:00:00 2001 From: Sergei Vinogradov Date: Wed, 20 Aug 2025 14:37:31 +0200 Subject: [PATCH 1/6] Move Kernel specific data from handler_impl to a separate data structure to use it in handler-based and handler-less submission paths --- sycl/include/sycl/handler.hpp | 11 ++- sycl/source/detail/handler_impl.hpp | 12 +-- sycl/source/detail/kernel_data.hpp | 81 +++++++++++++++++++ sycl/source/detail/queue_impl.hpp | 2 +- sycl/source/handler.cpp | 66 ++++++++------- sycl/test/abi/sycl_symbols_linux.dump | 1 - sycl/test/abi/sycl_symbols_windows.dump | 1 - .../arg_mask/EliminatedArgMask.cpp | 2 +- .../scheduler/SchedulerTestUtils.hpp | 8 +- .../scheduler/StreamInitDependencyOnHost.cpp | 7 +- 10 files changed, 133 insertions(+), 58 deletions(-) create mode 100644 sycl/source/detail/kernel_data.hpp diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index a17358775e812..996d3fbaf7f7a 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -856,14 +856,11 @@ class __SYCL_EXPORT handler { // Force hasSpecialCaptures to be evaluated at compile-time. constexpr bool HasSpecialCapt = detail::hasSpecialCaptures(); setKernelInfo((void *)MHostKernel->getPtr(), - detail::getKernelNumParams(), - &(detail::getKernelParamDesc), - detail::isKernelESIMD(), HasSpecialCapt); + &detail::getDeviceKernelInfo()); constexpr std::string_view KernelNameStr = detail::getKernelName(); MKernelName = KernelNameStr; - setDeviceKernelInfoPtr(&detail::getDeviceKernelInfo()); } else { // In case w/o the integration header it is necessary to process // accessors from the list(which are associated with this handler) as @@ -3658,10 +3655,13 @@ class __SYCL_EXPORT handler { void setNDRangeDescriptor(sycl::range<1> NumWorkItems, sycl::id<1> Offset); void setNDRangeDescriptor(sycl::range<1> NumWorkItems, sycl::range<1> LocalSize, sycl::id<1> Offset); - +#ifndef __INTEL_PREVIEW_BREAKING_CHANGES void setKernelInfo(void *KernelFuncPtr, int KernelNumArgs, detail::kernel_param_desc_t (*KernelParamDescGetter)(int), bool KernelIsESIMD, bool KernelHasSpecialCaptures); +#endif + void setKernelInfo(void *KernelFuncPtr, + detail::DeviceKernelInfo *DeviceKernelInfoPtr); void instantiateKernelOnHost(void *InstantiateKernelOnHostPtr); @@ -3689,7 +3689,6 @@ class __SYCL_EXPORT handler { void setKernelNameBasedCachePtr( detail::KernelNameBasedCacheT *KernelNameBasedCachePtr); #endif - void setDeviceKernelInfoPtr(detail::DeviceKernelInfo *DeviceKernelInfoPtr); queue getQueue(); diff --git a/sycl/source/detail/handler_impl.hpp b/sycl/source/detail/handler_impl.hpp index a8b217a1e64fa..a3cae92a58340 100644 --- a/sycl/source/detail/handler_impl.hpp +++ b/sycl/source/detail/handler_impl.hpp @@ -11,6 +11,7 @@ #include "sycl/handler.hpp" #include #include +#include #include #include @@ -236,16 +237,7 @@ class handler_impl { // Allocation ptr to be freed asynchronously. void *MFreePtr = nullptr; - // Store information about the kernel arguments. - void *MKernelFuncPtr = nullptr; - int MKernelNumArgs = 0; - detail::kernel_param_desc_t (*MKernelParamDescGetter)(int) = nullptr; - bool MKernelIsESIMD = false; - bool MKernelHasSpecialCaptures = true; - - // A pointer to device kernel information. Cached on the application side in - // headers or retrieved from program manager. - DeviceKernelInfo *MDeviceKernelInfoPtr = nullptr; + KernelData MKernelData; }; } // namespace detail diff --git a/sycl/source/detail/kernel_data.hpp b/sycl/source/detail/kernel_data.hpp new file mode 100644 index 0000000000000..d4ac09c8e383f --- /dev/null +++ b/sycl/source/detail/kernel_data.hpp @@ -0,0 +1,81 @@ +//==---------------- kernel_data.hpp - SYCL handler -----------------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#pragma once + +#include +#include + +namespace sycl { +inline namespace _V1 { +namespace detail { + +class KernelData { +public: + using KernelParamDescGetterT = detail::kernel_param_desc_t (*)(int); + + KernelData() = default; + ~KernelData() = default; + KernelData(const KernelData &) = default; + KernelData(KernelData &&) = default; + KernelData &operator=(const KernelData &) = default; + KernelData &operator=(KernelData &&) = default; + + void *getKernelFuncPtr() const { return MKernelFuncPtr; } + + size_t getKernelNumArgs() const { return MDeviceKernelInfoPtr->NumParams; } + + KernelParamDescGetterT getKernelParamDescGetter() const { + return MDeviceKernelInfoPtr->ParamDescGetter; + } + + bool isESIMD() const { return MDeviceKernelInfoPtr->IsESIMD; } + + bool hasSpecialCaptures() const { + return MDeviceKernelInfoPtr->HasSpecialCaptures; + } + + DeviceKernelInfo *getDeviceKernelInfoPtr() const { + return MDeviceKernelInfoPtr; + } + + void setDeviceKernelInfoPtr(DeviceKernelInfo *Ptr) { + MDeviceKernelInfoPtr = Ptr; + } +#ifndef __INTEL_PREVIEW_BREAKING_CHANGES + void setKernelInfo(void *KernelFuncPtr, int KernelNumArgs, + KernelParamDescGetterT KernelParamDescGetter, + bool KernelIsESIMD, bool KernelHasSpecialCaptures) { + MKernelFuncPtr = KernelFuncPtr; + MDeviceKernelInfoPtr->NumParams = KernelNumArgs; + MDeviceKernelInfoPtr->ParamDescGetter = KernelParamDescGetter; + MDeviceKernelInfoPtr->IsESIMD = KernelIsESIMD; + MDeviceKernelInfoPtr->HasSpecialCaptures = KernelHasSpecialCaptures; + } +#endif + + void setKernelInfo(void *KernelFuncPtr, + detail::DeviceKernelInfo *DeviceKernelInfoPtr) { + MKernelFuncPtr = KernelFuncPtr; + MDeviceKernelInfoPtr = DeviceKernelInfoPtr; + } + + bool usesAssert() const { return MDeviceKernelInfoPtr->usesAssert(); } + +private: + // Store information about the kernel arguments. + void *MKernelFuncPtr = nullptr; + + // A pointer to device kernel information. Cached on the application side in + // headers or retrieved from program manager. + DeviceKernelInfo *MDeviceKernelInfoPtr = nullptr; +}; + +} // namespace detail +} // namespace _V1 +} // namespace sycl diff --git a/sycl/source/detail/queue_impl.hpp b/sycl/source/detail/queue_impl.hpp index ca800e1511032..aca093debd622 100644 --- a/sycl/source/detail/queue_impl.hpp +++ b/sycl/source/detail/queue_impl.hpp @@ -872,7 +872,7 @@ class queue_impl : public std::enable_shared_from_this { // Kernel only uses assert if it's non interop one KernelUsesAssert = (!Handler.MKernel || Handler.MKernel->hasSYCLMetadata()) && - Handler.impl->MDeviceKernelInfoPtr->usesAssert(); + Handler.impl->MKernelData.usesAssert(); auto &PostProcess = *PostProcessorFunc; PostProcess(IsKernel, KernelUsesAssert, Event); diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index 397f01983add4..b05579f0945b4 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -496,12 +496,13 @@ event handler::finalize() { // Extract arguments from the kernel lambda, if required. // Skipping this is currently limited to simple kernels on the fast path. - if (type == detail::CGType::Kernel && impl->MKernelFuncPtr && - (!KernelFastPath || impl->MKernelHasSpecialCaptures)) { + if (type == detail::CGType::Kernel && impl->MKernelData.getKernelFuncPtr() && + (!KernelFastPath || impl->MKernelData.hasSpecialCaptures())) { clearArgs(); - extractArgsAndReqsFromLambda((char *)impl->MKernelFuncPtr, - impl->MKernelParamDescGetter, - impl->MKernelNumArgs, impl->MKernelIsESIMD); + extractArgsAndReqsFromLambda((char *)impl->MKernelData.getKernelFuncPtr(), + impl->MKernelData.getKernelParamDescGetter(), + impl->MKernelData.getKernelNumArgs(), + impl->MKernelData.isESIMD()); } // According to 4.7.6.9 of SYCL2020 spec, if a placeholder accessor is passed @@ -542,16 +543,17 @@ event handler::finalize() { } if (type == detail::CGType::Kernel) { - if (impl->MDeviceKernelInfoPtr) { + if (impl->MKernelData.getDeviceKernelInfoPtr()) { #ifndef __INTEL_PREVIEW_BREAKING_CHANGES - impl->MDeviceKernelInfoPtr->initIfNeeded(toKernelNameStrT(MKernelName)); + impl->MKernelData.getDeviceKernelInfoPtr()->initIfNeeded( + toKernelNameStrT(MKernelName)); #endif } else { // Fetch the device kernel info pointer if it hasn't been set (e.g. // in kernel bundle or free function cases). - impl->MDeviceKernelInfoPtr = + impl->MKernelData.setDeviceKernelInfoPtr( &detail::ProgramManager::getInstance().getOrCreateDeviceKernelInfo( - toKernelNameStrT(MKernelName)); + toKernelNameStrT(MKernelName))); } // If there were uses of set_specialization_constant build the kernel_bundle detail::kernel_bundle_impl *KernelBundleImpPtr = @@ -627,7 +629,7 @@ event handler::finalize() { if (DiscardEvent) { // Kernel only uses assert if it's non interop one bool KernelUsesAssert = !(MKernel && MKernel->isInterop()) && - impl->MDeviceKernelInfoPtr->usesAssert(); + impl->MKernelData.usesAssert(); DiscardEvent = !KernelUsesAssert; } @@ -647,7 +649,7 @@ event handler::finalize() { if (xptiEnabled) { std::tie(CmdTraceEvent, InstanceID) = emitKernelInstrumentationData( detail::GSYCLStreamID, MKernel, MCodeLoc, impl->MIsTopCodeLoc, - MKernelName.data(), *impl->MDeviceKernelInfoPtr, + MKernelName.data(), *impl->MKernelData.getDeviceKernelInfoPtr(), impl->get_queue_or_null(), impl->MNDRDesc, KernelBundleImpPtr, impl->MArgs); detail::emitInstrumentationGeneral(detail::GSYCLStreamID, InstanceID, @@ -664,11 +666,14 @@ event handler::finalize() { enqueueImpKernel( impl->get_queue(), impl->MNDRDesc, impl->MArgs, KernelBundleImpPtr, MKernel.get(), toKernelNameStrT(MKernelName), - *impl->MDeviceKernelInfoPtr, RawEvents, ResultEvent.get(), nullptr, - impl->MKernelCacheConfig, impl->MKernelIsCooperative, - impl->MKernelUsesClusterLaunch, impl->MKernelWorkGroupMemorySize, - BinImage, impl->MKernelFuncPtr, impl->MKernelNumArgs, - impl->MKernelParamDescGetter, impl->MKernelHasSpecialCaptures); + *impl->MKernelData.getDeviceKernelInfoPtr(), RawEvents, + ResultEvent.get(), nullptr, impl->MKernelCacheConfig, + impl->MKernelIsCooperative, impl->MKernelUsesClusterLaunch, + impl->MKernelWorkGroupMemorySize, BinImage, + impl->MKernelData.getKernelFuncPtr(), + impl->MKernelData.getKernelNumArgs(), + impl->MKernelData.getKernelParamDescGetter(), + impl->MKernelData.hasSpecialCaptures()); #ifdef XPTI_ENABLE_INSTRUMENTATION if (xptiEnabled) { // Emit signal only when event is created @@ -726,7 +731,7 @@ event handler::finalize() { impl->MNDRDesc, std::move(MHostKernel), std::move(MKernel), std::move(impl->MKernelBundle), std::move(impl->CGData), std::move(impl->MArgs), toKernelNameStrT(MKernelName), - *impl->MDeviceKernelInfoPtr, std::move(MStreamStorage), + *impl->MKernelData.getDeviceKernelInfoPtr(), std::move(MStreamStorage), std::move(impl->MAuxiliaryResources), getType(), impl->MKernelCacheConfig, impl->MKernelIsCooperative, impl->MKernelUsesClusterLaunch, impl->MKernelWorkGroupMemorySize, @@ -2610,26 +2615,25 @@ void handler::setNDRangeDescriptor(sycl::range<1> NumWorkItems, #ifndef __INTEL_PREVIEW_BREAKING_CHANGES void handler::setKernelNameBasedCachePtr( sycl::detail::KernelNameBasedCacheT *KernelNameBasedCachePtr) { - setDeviceKernelInfoPtr(reinterpret_cast( - KernelNameBasedCachePtr)); -} -#endif - -void handler::setDeviceKernelInfoPtr( - sycl::detail::DeviceKernelInfo *DeviceKernelInfoPtr) { - assert(!impl->MDeviceKernelInfoPtr && "Already set!"); - impl->MDeviceKernelInfoPtr = DeviceKernelInfoPtr; + assert(!impl->MKernelData.getDeviceKernelInfoPtr() && "Already set!"); + impl->MKernelData.setDeviceKernelInfoPtr( + reinterpret_cast( + KernelNameBasedCachePtr)); } void handler::setKernelInfo( void *KernelFuncPtr, int KernelNumArgs, detail::kernel_param_desc_t (*KernelParamDescGetter)(int), bool KernelIsESIMD, bool KernelHasSpecialCaptures) { - impl->MKernelFuncPtr = KernelFuncPtr; - impl->MKernelNumArgs = KernelNumArgs; - impl->MKernelParamDescGetter = KernelParamDescGetter; - impl->MKernelIsESIMD = KernelIsESIMD; - impl->MKernelHasSpecialCaptures = KernelHasSpecialCaptures; + impl->MKernelData.setKernelInfo(KernelFuncPtr, KernelNumArgs, + KernelParamDescGetter, KernelIsESIMD, + KernelHasSpecialCaptures); +} +#endif + +void handler::setKernelInfo(void *KernelFuncPtr, + detail::DeviceKernelInfo *DeviceKernelInfoPtr) { + impl->MKernelData.setKernelInfo(KernelFuncPtr, DeviceKernelInfoPtr); } void handler::instantiateKernelOnHost(void *InstantiateKernelOnHostPtr) { diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index e9f65ce662488..fd08eab77fdf3 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3592,7 +3592,6 @@ _ZN4sycl3_V17handler21setKernelWorkGroupMemEm _ZN4sycl3_V17handler21setUserFacingNodeTypeENS0_3ext6oneapi12experimental9node_typeE _ZN4sycl3_V17handler22ext_oneapi_fill2d_implEPvmPKvmmm _ZN4sycl3_V17handler22memcpyFromDeviceGlobalEPvPKvbmm -_ZN4sycl3_V17handler22setDeviceKernelInfoPtrEPNS0_6detail16DeviceKernelInfoE _ZN4sycl3_V17handler22setHandlerKernelBundleENS0_6kernelE _ZN4sycl3_V17handler22setHandlerKernelBundleERKSt10shared_ptrINS0_6detail18kernel_bundle_implEE _ZN4sycl3_V17handler22setKernelClusterLaunchENS0_5rangeILi1EEE diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index eae017c88eac8..c604b13f40d02 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -4409,7 +4409,6 @@ ?setArgsHelper@handler@_V1@sycl@@AEAAXH@Z ?setArgsToAssociatedAccessors@handler@_V1@sycl@@AEAAXXZ ?setDevice@HostProfilingInfo@detail@_V1@sycl@@QEAAXPEAVdevice_impl@234@@Z -?setDeviceKernelInfoPtr@handler@_V1@sycl@@AEAAXPEAVDeviceKernelInfo@detail@23@@Z ?setHandlerKernelBundle@handler@_V1@sycl@@AEAAXAEBV?$shared_ptr@Vkernel_bundle_impl@detail@_V1@sycl@@@std@@@Z ?setHandlerKernelBundle@handler@_V1@sycl@@AEAAXVkernel@23@@Z ?setKernelCacheConfig@handler@_V1@sycl@@AEAAXW4StableKernelCacheConfig@123@@Z diff --git a/sycl/unittests/program_manager/arg_mask/EliminatedArgMask.cpp b/sycl/unittests/program_manager/arg_mask/EliminatedArgMask.cpp index e6543927894a4..aea61760c1e70 100644 --- a/sycl/unittests/program_manager/arg_mask/EliminatedArgMask.cpp +++ b/sycl/unittests/program_manager/arg_mask/EliminatedArgMask.cpp @@ -148,7 +148,7 @@ class MockHandler : public sycl::handler { std::move(impl->MNDRDesc), std::move(CGH->MHostKernel), std::move(CGH->MKernel), std::move(impl->MKernelBundle), std::move(impl->CGData), std::move(impl->MArgs), - CGH->MKernelName.data(), *impl->MDeviceKernelInfoPtr, + CGH->MKernelName.data(), *impl->MKernelData.getDeviceKernelInfo(), std::move(CGH->MStreamStorage), std::move(impl->MAuxiliaryResources), impl->MCGType, {}, impl->MKernelIsCooperative, impl->MKernelUsesClusterLaunch, impl->MKernelWorkGroupMemorySize, diff --git a/sycl/unittests/scheduler/SchedulerTestUtils.hpp b/sycl/unittests/scheduler/SchedulerTestUtils.hpp index c3bdb342170de..067d3d201e2dd 100644 --- a/sycl/unittests/scheduler/SchedulerTestUtils.hpp +++ b/sycl/unittests/scheduler/SchedulerTestUtils.hpp @@ -303,10 +303,10 @@ class MockHandlerCustomFinalize : public MockHandler { CommandGroup.reset(new sycl::detail::CGExecKernel( getNDRDesc(), std::move(getHostKernel()), getKernel(), std::move(impl->MKernelBundle), std::move(CGData), getArgs(), - getKernelName(), *impl->MDeviceKernelInfoPtr, getStreamStorage(), - impl->MAuxiliaryResources, getType(), {}, impl->MKernelIsCooperative, - impl->MKernelUsesClusterLaunch, impl->MKernelWorkGroupMemorySize, - getCodeLoc())); + getKernelName(), *impl->MKernelData.getDeviceKernelInfoPtr(), + getStreamStorage(), impl->MAuxiliaryResources, getType(), {}, + impl->MKernelIsCooperative, impl->MKernelUsesClusterLaunch, + impl->MKernelWorkGroupMemorySize, getCodeLoc())); break; } case sycl::detail::CGType::CodeplayHostTask: { diff --git a/sycl/unittests/scheduler/StreamInitDependencyOnHost.cpp b/sycl/unittests/scheduler/StreamInitDependencyOnHost.cpp index 2b6d0cc563431..7be59dcd1a538 100644 --- a/sycl/unittests/scheduler/StreamInitDependencyOnHost.cpp +++ b/sycl/unittests/scheduler/StreamInitDependencyOnHost.cpp @@ -33,9 +33,10 @@ class MockHandlerStreamInit : public MockHandler { detail::CG::StorageInitHelper(getArgsStorage(), getAccStorage(), getSharedPtrStorage(), getRequirements(), getEvents()), - getArgs(), getKernelName(), *impl->MDeviceKernelInfoPtr, - getStreamStorage(), std::move(impl->MAuxiliaryResources), getType(), - {}, impl->MKernelIsCooperative, impl->MKernelUsesClusterLaunch, + getArgs(), getKernelName(), + *impl->MKernelData.getDeviceKernelInfoPtr(), getStreamStorage(), + std::move(impl->MAuxiliaryResources), getType(), {}, + impl->MKernelIsCooperative, impl->MKernelUsesClusterLaunch, impl->MKernelWorkGroupMemorySize, getCodeLoc())); break; } From 11c23d2b6769320fdbe714cda68da55534cbefba Mon Sep 17 00:00:00 2001 From: Sergei Vinogradov Date: Wed, 20 Aug 2025 22:44:32 +0200 Subject: [PATCH 2/6] Move MArgs and MDynamicParameters to KernelData --- .../oneapi/experimental/work_group_memory.hpp | 5 +- sycl/include/sycl/handler.hpp | 13 +- sycl/include/sycl/stream.hpp | 3 +- sycl/source/CMakeLists.txt | 1 + sycl/source/detail/graph/dynamic_impl.cpp | 2 +- sycl/source/detail/graph/graph_impl.cpp | 2 +- sycl/source/detail/handler_impl.hpp | 9 - sycl/source/detail/kernel_data.cpp | 352 +++++++++++++++++ sycl/source/detail/kernel_data.hpp | 51 +++ sycl/source/handler.cpp | 363 ++---------------- .../arg_mask/EliminatedArgMask.cpp | 2 +- .../scheduler/SchedulerTestUtils.hpp | 4 +- 12 files changed, 450 insertions(+), 357 deletions(-) create mode 100644 sycl/source/detail/kernel_data.cpp diff --git a/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp b/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp index ebee7791b9841..099e2c92a2c4f 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp @@ -18,9 +18,10 @@ namespace sycl { inline namespace _V1 { -class handler; namespace detail { +class KernelData; + template struct is_unbounded_array : std::false_type {}; template struct is_unbounded_array : std::true_type {}; @@ -38,7 +39,7 @@ class work_group_memory_impl { private: size_t buffer_size; - friend class sycl::handler; + friend class KernelData; }; } // namespace detail diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 996d3fbaf7f7a..3d81825ab831b 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -501,23 +501,21 @@ class __SYCL_EXPORT handler { extractArgsAndReqsFromLambda(char *LambdaPtr, size_t KernelArgsNum, const detail::kernel_param_desc_t *KernelArgs, bool IsESIMD); -#endif /// Extracts and prepares kernel arguments from the lambda using information /// from the built-ins or integration header. void extractArgsAndReqsFromLambda( char *LambdaPtr, detail::kernel_param_desc_t (*ParamDescGetter)(int), size_t NumKernelParams, bool IsESIMD); - +#endif /// Extracts and prepares kernel arguments set via set_arg(s). void extractArgsAndReqs(); -#if defined(__INTEL_PREVIEW_BREAKING_CHANGES) - // TODO: processArg need not to be public - __SYCL_DLL_LOCAL -#endif +#ifndef __INTEL_PREVIEW_BREAKING_CHANGES + // TODO: remove in the next ABI-breaking window. void processArg(void *Ptr, const detail::kernel_param_kind_t &Kind, const int Size, const size_t Index, size_t &IndexShift, bool IsKernelCreatedFromSource, bool IsESIMD); +#endif /// \return a string containing name of SYCL kernel. detail::ABINeutralKernelNameStrT getKernelName(); @@ -3608,7 +3606,10 @@ class __SYCL_EXPORT handler { void addArg(detail::kernel_param_kind_t ArgKind, void *Req, int AccessTarget, int ArgIndex); +#ifndef __INTEL_PREVIEW_BREAKING_CHANGES + // TODO: remove in the next ABI-breaking window void clearArgs(); +#endif void setArgsToAssociatedAccessors(); bool HasAssociatedAccessor(detail::AccessorImplHost *Req, diff --git a/sycl/include/sycl/stream.hpp b/sycl/include/sycl/stream.hpp index b583b03635e82..560ca1b623126 100644 --- a/sycl/include/sycl/stream.hpp +++ b/sycl/include/sycl/stream.hpp @@ -42,6 +42,7 @@ inline namespace _V1 { namespace detail { class stream_impl; +class KernelData; using FmtFlags = unsigned int; @@ -1041,7 +1042,7 @@ class __SYCL_EXPORT __SYCL_SPECIAL_CLASS __SYCL_TYPE(stream) stream } #endif - friend class handler; + friend class detail::KernelData; template friend class ext::oneapi::weak_object; diff --git a/sycl/source/CMakeLists.txt b/sycl/source/CMakeLists.txt index 38ffd232fcbbe..fdc4fcd77bcac 100644 --- a/sycl/source/CMakeLists.txt +++ b/sycl/source/CMakeLists.txt @@ -267,6 +267,7 @@ set(SYCL_COMMON_SOURCES "detail/device_filter.cpp" "detail/host_pipe_map.cpp" "detail/device_global_map.cpp" + "detail/kernel_data.cpp" "detail/kernel_global_info.cpp" "detail/device_global_map_entry.cpp" "detail/device_image_impl.cpp" diff --git a/sycl/source/detail/graph/dynamic_impl.cpp b/sycl/source/detail/graph/dynamic_impl.cpp index 3944dc6a2cfc3..c04d424b84dbf 100644 --- a/sycl/source/detail/graph/dynamic_impl.cpp +++ b/sycl/source/detail/graph/dynamic_impl.cpp @@ -343,7 +343,7 @@ void dynamic_command_group_impl::finalizeCGFList( MCommandGroups.push_back(std::shared_ptr(RawCGPtr)); // Track dynamic_parameter usage in command-group - auto &DynamicParams = Handler.impl->MDynamicParameters; + auto &DynamicParams = Handler.impl->MKernelData.getDynamicParameters(); if (DynamicParams.size() > 0 && Handler.getType() == sycl::detail::CGType::CodeplayHostTask) { diff --git a/sycl/source/detail/graph/graph_impl.cpp b/sycl/source/detail/graph/graph_impl.cpp index 4583efbe881dc..ed4ef87cd0916 100644 --- a/sycl/source/detail/graph/graph_impl.cpp +++ b/sycl/source/detail/graph/graph_impl.cpp @@ -471,7 +471,7 @@ node_impl &graph_impl::add(std::function CGF, // Retrieve any dynamic parameters which have been registered in the CGF and // register the actual nodes with them. - auto &DynamicParams = Handler.impl->MDynamicParameters; + auto &DynamicParams = Handler.impl->MKernelData.getDynamicParameters(); if (NodeType != node_type::kernel && DynamicParams.size() > 0) { throw sycl::exception(sycl::make_error_code(errc::invalid), diff --git a/sycl/source/detail/handler_impl.hpp b/sycl/source/detail/handler_impl.hpp index a3cae92a58340..6235f49f26697 100644 --- a/sycl/source/detail/handler_impl.hpp +++ b/sycl/source/detail/handler_impl.hpp @@ -134,21 +134,12 @@ class handler_impl { sycl::ext::oneapi::experimental::node_type MUserFacingNodeType = sycl::ext::oneapi::experimental::node_type::empty; - // Storage for any SYCL Graph dynamic parameters which have been flagged for - // registration in the CG, along with the argument index for the parameter. - std::vector> - MDynamicParameters; - /// The storage for the arguments passed. /// We need to store a copy of values that are passed explicitly through /// set_arg, require and so on, because we need them to be alive after /// we exit the method they are passed in. detail::CG::StorageInitHelper CGData; - /// The list of arguments for the kernel. - std::vector MArgs; - /// The list of associated accessors with this handler. /// These accessors were created with this handler as argument or /// have become required for this handler via require method. diff --git a/sycl/source/detail/kernel_data.cpp b/sycl/source/detail/kernel_data.cpp new file mode 100644 index 0000000000000..4771b67453d15 --- /dev/null +++ b/sycl/source/detail/kernel_data.cpp @@ -0,0 +1,352 @@ +//==-------------------- kernel_data.cpp ----------------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include + +#include + +#include +#include + +#include + +namespace sycl { +inline namespace _V1 { +namespace detail { + +// The argument can take up more space to store additional information about +// MAccessRange, MMemoryRange, and MOffset added with addArgsForGlobalAccessor. +// We use the worst-case estimate because the lifetime of the vector is short. +// In processArg the kind_stream case introduces the maximum number of +// additional arguments. The case adds additional 12 arguments to the currently +// processed argument, hence worst-case estimate is 12+1=13. +// TODO: the constant can be removed if the size of MArgs will be calculated at +// compile time. +inline constexpr size_t MaxNumAdditionalArgs = 13; + +constexpr static int AccessTargetMask = 0x7ff; + +static void addArgsForGlobalAccessor(detail::Requirement *AccImpl, size_t Index, + size_t &IndexShift, int Size, + bool IsKernelCreatedFromSource, + size_t GlobalSize, + std::vector &Args, + bool isESIMD) { + using detail::kernel_param_kind_t; + if (AccImpl->PerWI) + AccImpl->resize(GlobalSize); + + Args.emplace_back(kernel_param_kind_t::kind_accessor, AccImpl, Size, + Index + IndexShift); + + // TODO ESIMD currently does not suport offset, memory and access ranges - + // accessor::init for ESIMD-mode accessor has a single field, translated + // to a single kernel argument set above. + if (!isESIMD && !IsKernelCreatedFromSource) { + // Dimensionality of the buffer is 1 when dimensionality of the + // accessor is 0. + const size_t SizeAccField = + sizeof(size_t) * (AccImpl->MDims == 0 ? 1 : AccImpl->MDims); + ++IndexShift; + Args.emplace_back(kernel_param_kind_t::kind_std_layout, + &AccImpl->MAccessRange[0], SizeAccField, + Index + IndexShift); + ++IndexShift; + Args.emplace_back(kernel_param_kind_t::kind_std_layout, + &AccImpl->MMemoryRange[0], SizeAccField, + Index + IndexShift); + ++IndexShift; + Args.emplace_back(kernel_param_kind_t::kind_std_layout, + &AccImpl->MOffset[0], SizeAccField, Index + IndexShift); + } +} + +static void addArgsForLocalAccessor(detail::LocalAccessorImplHost *LAcc, + size_t Index, size_t &IndexShift, + bool IsKernelCreatedFromSource, + std::vector &Args, + bool IsESIMD) { + using detail::kernel_param_kind_t; + + range<3> &LAccSize = LAcc->MSize; + const int Dims = LAcc->MDims; + int SizeInBytes = LAcc->MElemSize; + for (int I = 0; I < Dims; ++I) + SizeInBytes *= LAccSize[I]; + + // Some backends do not accept zero-sized local memory arguments, so we + // make it a minimum allocation of 1 byte. + SizeInBytes = std::max(SizeInBytes, 1); + Args.emplace_back(kernel_param_kind_t::kind_std_layout, nullptr, SizeInBytes, + Index + IndexShift); + // TODO ESIMD currently does not suport MSize field passing yet + // accessor::init for ESIMD-mode accessor has a single field, translated + // to a single kernel argument set above. + if (!IsESIMD && !IsKernelCreatedFromSource) { + ++IndexShift; + const size_t SizeAccField = (Dims == 0 ? 1 : Dims) * sizeof(LAccSize[0]); + Args.emplace_back(kernel_param_kind_t::kind_std_layout, &LAccSize, + SizeAccField, Index + IndexShift); + ++IndexShift; + Args.emplace_back(kernel_param_kind_t::kind_std_layout, &LAccSize, + SizeAccField, Index + IndexShift); + ++IndexShift; + Args.emplace_back(kernel_param_kind_t::kind_std_layout, &LAccSize, + SizeAccField, Index + IndexShift); + } +} + +void KernelData::processArg(void *Ptr, const detail::kernel_param_kind_t &Kind, + const int Size, const size_t Index, + size_t &IndexShift, bool IsKernelCreatedFromSource, + const NDRDescT &NDRDesc) { + using detail::kernel_param_kind_t; + size_t GlobalSize = NDRDesc.GlobalSize[0]; + for (size_t I = 1; I < NDRDesc.Dims; ++I) { + GlobalSize *= NDRDesc.GlobalSize[I]; + } + + switch (Kind) { + case kernel_param_kind_t::kind_std_layout: + case kernel_param_kind_t::kind_pointer: { + addArg(Kind, Ptr, Size, Index + IndexShift); + break; + } + case kernel_param_kind_t::kind_stream: { + // Stream contains several accessors inside. + stream *S = static_cast(Ptr); + + detail::AccessorBaseHost *GBufBase = + static_cast(&S->GlobalBuf); + detail::Requirement *GBufReq = &*detail::getSyclObjImpl(*GBufBase); + addArgsForGlobalAccessor(GBufReq, Index, IndexShift, Size, + IsKernelCreatedFromSource, GlobalSize, MArgs, + isESIMD()); + ++IndexShift; + detail::AccessorBaseHost *GOffsetBase = + static_cast(&S->GlobalOffset); + detail::Requirement *GOffsetReq = &*detail::getSyclObjImpl(*GOffsetBase); + addArgsForGlobalAccessor(GOffsetReq, Index, IndexShift, Size, + IsKernelCreatedFromSource, GlobalSize, MArgs, + isESIMD()); + ++IndexShift; + detail::AccessorBaseHost *GFlushBase = + static_cast(&S->GlobalFlushBuf); + detail::Requirement *GFlushReq = &*detail::getSyclObjImpl(*GFlushBase); + + // If work group size wasn't set explicitly then it must be recieved + // from kernel attribute or set to default values. + // For now we can't get this attribute here. + // So we just suppose that WG size is always default for stream. + // TODO adjust MNDRDesc when device image contains kernel's attribute + if (GlobalSize == 0) { + GlobalSize = NDRDesc.NumWorkGroups[0]; + for (size_t I = 1; I < NDRDesc.Dims; ++I) { + GlobalSize *= NDRDesc.NumWorkGroups[I]; + } + } + addArgsForGlobalAccessor(GFlushReq, Index, IndexShift, Size, + IsKernelCreatedFromSource, GlobalSize, MArgs, + isESIMD()); + ++IndexShift; + addArg(kernel_param_kind_t::kind_std_layout, &S->FlushBufferSize, + sizeof(S->FlushBufferSize), Index + IndexShift); + + break; + } + case kernel_param_kind_t::kind_accessor: { + // For args kind of accessor Size is information about accessor. + // The first 11 bits of Size encodes the accessor target. + const access::target AccTarget = + static_cast(Size & AccessTargetMask); + switch (AccTarget) { + case access::target::device: + case access::target::constant_buffer: { + detail::Requirement *AccImpl = static_cast(Ptr); + addArgsForGlobalAccessor(AccImpl, Index, IndexShift, Size, + IsKernelCreatedFromSource, GlobalSize, MArgs, + isESIMD()); + break; + } + case access::target::local: { + detail::LocalAccessorImplHost *LAccImpl = + static_cast(Ptr); + + addArgsForLocalAccessor(LAccImpl, Index, IndexShift, + IsKernelCreatedFromSource, MArgs, isESIMD()); + break; + } + case access::target::image: + case access::target::image_array: { + detail::Requirement *AccImpl = static_cast(Ptr); + addArg(Kind, AccImpl, Size, Index + IndexShift); + if (!IsKernelCreatedFromSource) { + // TODO Handle additional kernel arguments for image class + // if the compiler front-end adds them. + } + break; + } + case access::target::host_image: + case access::target::host_task: + case access::target::host_buffer: { + throw sycl::exception(make_error_code(errc::invalid), + "Unsupported accessor target case."); + break; + } + } + break; + } + case kernel_param_kind_t::kind_dynamic_accessor: { + const access::target AccTarget = + static_cast(Size & AccessTargetMask); + switch (AccTarget) { + case access::target::local: { + + // We need to recover the inheritance layout by casting to + // dynamic_parameter_impl first. Casting directly to + // dynamic_local_accessor_impl would result in an incorrect pointer. + auto *DynParamImpl = static_cast< + ext::oneapi::experimental::detail::dynamic_parameter_impl *>(Ptr); + + MDynamicParameters.emplace_back(DynParamImpl, Index + IndexShift); + + auto *DynLocalAccessorImpl = static_cast< + ext::oneapi::experimental::detail::dynamic_local_accessor_impl *>( + DynParamImpl); + + addArgsForLocalAccessor(&DynLocalAccessorImpl->LAccImplHost, Index, + IndexShift, IsKernelCreatedFromSource, MArgs, + isESIMD()); + break; + } + default: { + assert(false && "Unsupported dynamic accessor target"); + } + } + break; + } + case kernel_param_kind_t::kind_dynamic_work_group_memory: { + + // We need to recover the inheritance layout by casting to + // dynamic_parameter_impl first. Casting directly to + // dynamic_work_group_memory_impl would result in an incorrect pointer. + auto *DynParamImpl = static_cast< + ext::oneapi::experimental::detail::dynamic_parameter_impl *>(Ptr); + + MDynamicParameters.emplace_back(DynParamImpl, Index + IndexShift); + + auto *DynWorkGroupImpl = static_cast< + ext::oneapi::experimental::detail::dynamic_work_group_memory_impl *>( + DynParamImpl); + + addArg(kernel_param_kind_t::kind_std_layout, nullptr, + DynWorkGroupImpl->BufferSizeInBytes, Index + IndexShift); + break; + } + case kernel_param_kind_t::kind_work_group_memory: { + addArg(kernel_param_kind_t::kind_std_layout, nullptr, + static_cast(Ptr)->buffer_size, + Index + IndexShift); + break; + } + case kernel_param_kind_t::kind_sampler: { + addArg(kernel_param_kind_t::kind_sampler, Ptr, sizeof(sampler), + Index + IndexShift); + break; + } + case kernel_param_kind_t::kind_specialization_constants_buffer: { + addArg(kernel_param_kind_t::kind_specialization_constants_buffer, Ptr, Size, + Index + IndexShift); + break; + } + case kernel_param_kind_t::kind_invalid: + throw exception(make_error_code(errc::invalid), + "Invalid kernel param kind"); + break; + } +} + +void KernelData::extractArgsAndReqs(const NDRDescT &NDRDesc, + bool IsKernelCreatedFromSource) { + std::vector UnPreparedArgs = std::move(MArgs); + clearArgs(); + + std::sort( + UnPreparedArgs.begin(), UnPreparedArgs.end(), + [](const detail::ArgDesc &first, const detail::ArgDesc &second) -> bool { + return (first.MIndex < second.MIndex); + }); + + MArgs.reserve(MaxNumAdditionalArgs * UnPreparedArgs.size()); + + size_t IndexShift = 0; + for (size_t I = 0; I < UnPreparedArgs.size(); ++I) { + void *Ptr = UnPreparedArgs[I].MPtr; + const detail::kernel_param_kind_t &Kind = UnPreparedArgs[I].MType; + const int &Size = UnPreparedArgs[I].MSize; + const int Index = UnPreparedArgs[I].MIndex; + processArg(Ptr, Kind, Size, Index, IndexShift, IsKernelCreatedFromSource, + NDRDesc); + } +} + +void KernelData::extractArgsAndReqsFromLambda(const NDRDescT &NDRDesc) { + size_t IndexShift = 0; + clearArgs(); + MArgs.reserve(MaxNumAdditionalArgs * getKernelNumArgs()); + + for (size_t I = 0; I < getKernelNumArgs(); ++I) { + auto KernelParamDescGetter = getKernelParamDescGetter(); + detail::kernel_param_desc_t ParamDesc = KernelParamDescGetter(I); + void *Ptr = (char *)MKernelFuncPtr + ParamDesc.offset; + const detail::kernel_param_kind_t &Kind = ParamDesc.kind; + const int &Size = ParamDesc.info; + if (Kind == detail::kernel_param_kind_t::kind_accessor) { + // For args kind of accessor Size is information about accessor. + // The first 11 bits of Size encodes the accessor target. + const access::target AccTarget = + static_cast(Size & AccessTargetMask); + if ((AccTarget == access::target::device || + AccTarget == access::target::constant_buffer) || + (AccTarget == access::target::image || + AccTarget == access::target::image_array)) { + detail::AccessorBaseHost *AccBase = + static_cast(Ptr); + Ptr = detail::getSyclObjImpl(*AccBase).get(); + } else if (AccTarget == access::target::local) { + detail::LocalAccessorBaseHost *LocalAccBase = + static_cast(Ptr); + Ptr = detail::getSyclObjImpl(*LocalAccBase).get(); + } + } else if (Kind == detail::kernel_param_kind_t::kind_dynamic_accessor) { + // For args kind of accessor Size is information about accessor. + // The first 11 bits of Size encodes the accessor target. + // Only local targets are supported for dynamic accessors. + assert(static_cast(Size & AccessTargetMask) == + access::target::local); + + ext::oneapi::experimental::detail::dynamic_parameter_base + *DynamicParamBase = static_cast< + ext::oneapi::experimental::detail::dynamic_parameter_base *>(Ptr); + Ptr = detail::getSyclObjImpl(*DynamicParamBase).get(); + } else if (Kind == + detail::kernel_param_kind_t::kind_dynamic_work_group_memory) { + ext::oneapi::experimental::detail::dynamic_parameter_base + *DynamicParamBase = static_cast< + ext::oneapi::experimental::detail::dynamic_parameter_base *>(Ptr); + Ptr = detail::getSyclObjImpl(*DynamicParamBase).get(); + } + + processArg(Ptr, Kind, Size, I, IndexShift, + /*IsKernelCreatedFromSource=*/false, NDRDesc); + } +} + +} // namespace detail +} // namespace _V1 +} // namespace sycl diff --git a/sycl/source/detail/kernel_data.hpp b/sycl/source/detail/kernel_data.hpp index d4ac09c8e383f..f1891347cc9ab 100644 --- a/sycl/source/detail/kernel_data.hpp +++ b/sycl/source/detail/kernel_data.hpp @@ -8,9 +8,15 @@ #pragma once +#include +#include +#include + #include #include +#include + namespace sycl { inline namespace _V1 { namespace detail { @@ -18,6 +24,9 @@ namespace detail { class KernelData { public: using KernelParamDescGetterT = detail::kernel_param_desc_t (*)(int); + using DynamicParametersVecT = std::vector>; + using ArgsVecT = std::vector; KernelData() = default; ~KernelData() = default; @@ -26,6 +35,32 @@ class KernelData { KernelData &operator=(const KernelData &) = default; KernelData &operator=(KernelData &&) = default; + DynamicParametersVecT &getDynamicParameters() { return MDynamicParameters; } + + const DynamicParametersVecT &getDynamicParameters() const { + return MDynamicParameters; + } + + template void addDynamicParameter(Args &&...args) { + MDynamicParameters.emplace_back(std::forward(args)...); + } + + ArgsVecT &getArgs() & { return MArgs; } + + const ArgsVecT &getArgs() const & { return MArgs; } + + ArgsVecT &&getArgs() && { return std::move(MArgs); } + + void setArgs(const ArgsVecT &Args) { MArgs = Args; } + + void addArg(const detail::ArgDesc &Arg) { MArgs.push_back(Arg); } + + template void addArg(Args &&...args) { + MArgs.emplace_back(std::forward(args)...); + } + + void clearArgs() { MArgs.clear(); } + void *getKernelFuncPtr() const { return MKernelFuncPtr; } size_t getKernelNumArgs() const { return MDeviceKernelInfoPtr->NumParams; } @@ -67,7 +102,23 @@ class KernelData { bool usesAssert() const { return MDeviceKernelInfoPtr->usesAssert(); } + void processArg(void *Ptr, const detail::kernel_param_kind_t &Kind, + const int Size, const size_t Index, size_t &IndexShift, + bool IsKernelCreatedFromSource, const NDRDescT &NDRDesc); + + void extractArgsAndReqs(const NDRDescT &NDRDesc, + bool IsKernelCreatedFromSource); + + void extractArgsAndReqsFromLambda(const NDRDescT &NDRDesc); + private: + // Storage for any SYCL Graph dynamic parameters which have been flagged for + // registration in the CG, along with the argument index for the parameter. + DynamicParametersVecT MDynamicParameters; + + /// The list of arguments for the kernel. + std::vector MArgs; + // Store information about the kernel arguments. void *MKernelFuncPtr = nullptr; diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index b05579f0945b4..2986b30120da8 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -498,18 +498,14 @@ event handler::finalize() { // Skipping this is currently limited to simple kernels on the fast path. if (type == detail::CGType::Kernel && impl->MKernelData.getKernelFuncPtr() && (!KernelFastPath || impl->MKernelData.hasSpecialCaptures())) { - clearArgs(); - extractArgsAndReqsFromLambda((char *)impl->MKernelData.getKernelFuncPtr(), - impl->MKernelData.getKernelParamDescGetter(), - impl->MKernelData.getKernelNumArgs(), - impl->MKernelData.isESIMD()); + impl->MKernelData.extractArgsAndReqsFromLambda(impl->MNDRDesc); } // According to 4.7.6.9 of SYCL2020 spec, if a placeholder accessor is passed // to a command without being bound to a command group, an exception should // be thrown. { - for (const auto &arg : impl->MArgs) { + for (const auto &arg : impl->MKernelData.getArgs()) { if (arg.MType != detail::kernel_param_kind_t::kind_accessor) continue; @@ -651,7 +647,7 @@ event handler::finalize() { detail::GSYCLStreamID, MKernel, MCodeLoc, impl->MIsTopCodeLoc, MKernelName.data(), *impl->MKernelData.getDeviceKernelInfoPtr(), impl->get_queue_or_null(), impl->MNDRDesc, KernelBundleImpPtr, - impl->MArgs); + impl->MKernelData.getArgs()); detail::emitInstrumentationGeneral(detail::GSYCLStreamID, InstanceID, CmdTraceEvent, xpti::trace_task_begin, nullptr); @@ -664,8 +660,8 @@ event handler::finalize() { assert(BinImage && "Failed to obtain a binary image."); } enqueueImpKernel( - impl->get_queue(), impl->MNDRDesc, impl->MArgs, KernelBundleImpPtr, - MKernel.get(), toKernelNameStrT(MKernelName), + impl->get_queue(), impl->MNDRDesc, impl->MKernelData.getArgs(), + KernelBundleImpPtr, MKernel.get(), toKernelNameStrT(MKernelName), *impl->MKernelData.getDeviceKernelInfoPtr(), RawEvents, ResultEvent.get(), nullptr, impl->MKernelCacheConfig, impl->MKernelIsCooperative, impl->MKernelUsesClusterLaunch, @@ -730,7 +726,7 @@ event handler::finalize() { CommandGroup.reset(new detail::CGExecKernel( impl->MNDRDesc, std::move(MHostKernel), std::move(MKernel), std::move(impl->MKernelBundle), std::move(impl->CGData), - std::move(impl->MArgs), toKernelNameStrT(MKernelName), + std::move(impl->MKernelData).getArgs(), toKernelNameStrT(MKernelName), *impl->MKernelData.getDeviceKernelInfoPtr(), std::move(MStreamStorage), std::move(impl->MAuxiliaryResources), getType(), impl->MKernelCacheConfig, impl->MKernelIsCooperative, @@ -791,9 +787,10 @@ event handler::finalize() { case detail::CGType::CodeplayHostTask: { detail::context_impl &Context = impl->get_context(); detail::queue_impl *Queue = impl->get_queue_or_null(); - CommandGroup.reset(new detail::CGHostTask( - std::move(impl->MHostTask), Queue, &Context, std::move(impl->MArgs), - std::move(impl->CGData), getType(), MCodeLoc)); + CommandGroup.reset( + new detail::CGHostTask(std::move(impl->MHostTask), Queue, &Context, + std::move(impl->MKernelData).getArgs(), + std::move(impl->CGData), getType(), MCodeLoc)); break; } case detail::CGType::Barrier: @@ -1057,244 +1054,14 @@ void handler::associateWithHandler( static_cast(AccTarget)); } -static void addArgsForGlobalAccessor(detail::Requirement *AccImpl, size_t Index, - size_t &IndexShift, int Size, - bool IsKernelCreatedFromSource, - size_t GlobalSize, - std::vector &Args, - bool isESIMD) { - using detail::kernel_param_kind_t; - if (AccImpl->PerWI) - AccImpl->resize(GlobalSize); - - Args.emplace_back(kernel_param_kind_t::kind_accessor, AccImpl, Size, - Index + IndexShift); - - // TODO ESIMD currently does not suport offset, memory and access ranges - - // accessor::init for ESIMD-mode accessor has a single field, translated - // to a single kernel argument set above. - if (!isESIMD && !IsKernelCreatedFromSource) { - // Dimensionality of the buffer is 1 when dimensionality of the - // accessor is 0. - const size_t SizeAccField = - sizeof(size_t) * (AccImpl->MDims == 0 ? 1 : AccImpl->MDims); - ++IndexShift; - Args.emplace_back(kernel_param_kind_t::kind_std_layout, - &AccImpl->MAccessRange[0], SizeAccField, - Index + IndexShift); - ++IndexShift; - Args.emplace_back(kernel_param_kind_t::kind_std_layout, - &AccImpl->MMemoryRange[0], SizeAccField, - Index + IndexShift); - ++IndexShift; - Args.emplace_back(kernel_param_kind_t::kind_std_layout, - &AccImpl->MOffset[0], SizeAccField, Index + IndexShift); - } -} - -static void addArgsForLocalAccessor(detail::LocalAccessorImplHost *LAcc, - size_t Index, size_t &IndexShift, - bool IsKernelCreatedFromSource, - std::vector &Args, - bool IsESIMD) { - using detail::kernel_param_kind_t; - - range<3> &LAccSize = LAcc->MSize; - const int Dims = LAcc->MDims; - int SizeInBytes = LAcc->MElemSize; - for (int I = 0; I < Dims; ++I) - SizeInBytes *= LAccSize[I]; - - // Some backends do not accept zero-sized local memory arguments, so we - // make it a minimum allocation of 1 byte. - SizeInBytes = std::max(SizeInBytes, 1); - Args.emplace_back(kernel_param_kind_t::kind_std_layout, nullptr, SizeInBytes, - Index + IndexShift); - // TODO ESIMD currently does not suport MSize field passing yet - // accessor::init for ESIMD-mode accessor has a single field, translated - // to a single kernel argument set above. - if (!IsESIMD && !IsKernelCreatedFromSource) { - ++IndexShift; - const size_t SizeAccField = (Dims == 0 ? 1 : Dims) * sizeof(LAccSize[0]); - Args.emplace_back(kernel_param_kind_t::kind_std_layout, &LAccSize, - SizeAccField, Index + IndexShift); - ++IndexShift; - Args.emplace_back(kernel_param_kind_t::kind_std_layout, &LAccSize, - SizeAccField, Index + IndexShift); - ++IndexShift; - Args.emplace_back(kernel_param_kind_t::kind_std_layout, &LAccSize, - SizeAccField, Index + IndexShift); - } -} - +#ifndef __INTEL_PREVIEW_BREAKING_CHANGES void handler::processArg(void *Ptr, const detail::kernel_param_kind_t &Kind, const int Size, const size_t Index, size_t &IndexShift, bool IsKernelCreatedFromSource, bool IsESIMD) { - using detail::kernel_param_kind_t; - size_t GlobalSize = impl->MNDRDesc.GlobalSize[0]; - for (size_t I = 1; I < impl->MNDRDesc.Dims; ++I) { - GlobalSize *= impl->MNDRDesc.GlobalSize[I]; - } - - switch (Kind) { - case kernel_param_kind_t::kind_std_layout: - case kernel_param_kind_t::kind_pointer: { - addArg(Kind, Ptr, Size, Index + IndexShift); - break; - } - case kernel_param_kind_t::kind_stream: { - // Stream contains several accessors inside. - stream *S = static_cast(Ptr); - - detail::AccessorBaseHost *GBufBase = - static_cast(&S->GlobalBuf); - detail::Requirement *GBufReq = &*detail::getSyclObjImpl(*GBufBase); - addArgsForGlobalAccessor(GBufReq, Index, IndexShift, Size, - IsKernelCreatedFromSource, GlobalSize, impl->MArgs, - IsESIMD); - ++IndexShift; - detail::AccessorBaseHost *GOffsetBase = - static_cast(&S->GlobalOffset); - detail::Requirement *GOffsetReq = &*detail::getSyclObjImpl(*GOffsetBase); - addArgsForGlobalAccessor(GOffsetReq, Index, IndexShift, Size, - IsKernelCreatedFromSource, GlobalSize, impl->MArgs, - IsESIMD); - ++IndexShift; - detail::AccessorBaseHost *GFlushBase = - static_cast(&S->GlobalFlushBuf); - detail::Requirement *GFlushReq = &*detail::getSyclObjImpl(*GFlushBase); - - // If work group size wasn't set explicitly then it must be recieved - // from kernel attribute or set to default values. - // For now we can't get this attribute here. - // So we just suppose that WG size is always default for stream. - // TODO adjust MNDRDesc when device image contains kernel's attribute - if (GlobalSize == 0) { - GlobalSize = impl->MNDRDesc.NumWorkGroups[0]; - for (size_t I = 1; I < impl->MNDRDesc.Dims; ++I) { - GlobalSize *= impl->MNDRDesc.NumWorkGroups[I]; - } - } - addArgsForGlobalAccessor(GFlushReq, Index, IndexShift, Size, - IsKernelCreatedFromSource, GlobalSize, impl->MArgs, - IsESIMD); - ++IndexShift; - addArg(kernel_param_kind_t::kind_std_layout, &S->FlushBufferSize, - sizeof(S->FlushBufferSize), Index + IndexShift); - - break; - } - case kernel_param_kind_t::kind_accessor: { - // For args kind of accessor Size is information about accessor. - // The first 11 bits of Size encodes the accessor target. - const access::target AccTarget = - static_cast(Size & AccessTargetMask); - switch (AccTarget) { - case access::target::device: - case access::target::constant_buffer: { - detail::Requirement *AccImpl = static_cast(Ptr); - addArgsForGlobalAccessor(AccImpl, Index, IndexShift, Size, - IsKernelCreatedFromSource, GlobalSize, - impl->MArgs, IsESIMD); - break; - } - case access::target::local: { - detail::LocalAccessorImplHost *LAccImpl = - static_cast(Ptr); - - addArgsForLocalAccessor(LAccImpl, Index, IndexShift, - IsKernelCreatedFromSource, impl->MArgs, IsESIMD); - break; - } - case access::target::image: - case access::target::image_array: { - detail::Requirement *AccImpl = static_cast(Ptr); - addArg(Kind, AccImpl, Size, Index + IndexShift); - if (!IsKernelCreatedFromSource) { - // TODO Handle additional kernel arguments for image class - // if the compiler front-end adds them. - } - break; - } - case access::target::host_image: - case access::target::host_task: - case access::target::host_buffer: { - throw sycl::exception(make_error_code(errc::invalid), - "Unsupported accessor target case."); - break; - } - } - break; - } - case kernel_param_kind_t::kind_dynamic_accessor: { - const access::target AccTarget = - static_cast(Size & AccessTargetMask); - switch (AccTarget) { - case access::target::local: { - - // We need to recover the inheritance layout by casting to - // dynamic_parameter_impl first. Casting directly to - // dynamic_local_accessor_impl would result in an incorrect pointer. - auto *DynParamImpl = static_cast< - ext::oneapi::experimental::detail::dynamic_parameter_impl *>(Ptr); - - registerDynamicParameter(DynParamImpl, Index + IndexShift); - - auto *DynLocalAccessorImpl = static_cast< - ext::oneapi::experimental::detail::dynamic_local_accessor_impl *>( - DynParamImpl); - - addArgsForLocalAccessor(&DynLocalAccessorImpl->LAccImplHost, Index, - IndexShift, IsKernelCreatedFromSource, - impl->MArgs, IsESIMD); - break; - } - default: { - assert(false && "Unsupported dynamic accessor target"); - } - } - break; - } - case kernel_param_kind_t::kind_dynamic_work_group_memory: { - - // We need to recover the inheritance layout by casting to - // dynamic_parameter_impl first. Casting directly to - // dynamic_work_group_memory_impl would result in an incorrect pointer. - auto *DynParamImpl = static_cast< - ext::oneapi::experimental::detail::dynamic_parameter_impl *>(Ptr); - - registerDynamicParameter(DynParamImpl, Index + IndexShift); - - auto *DynWorkGroupImpl = static_cast< - ext::oneapi::experimental::detail::dynamic_work_group_memory_impl *>( - DynParamImpl); - - addArg(kernel_param_kind_t::kind_std_layout, nullptr, - DynWorkGroupImpl->BufferSizeInBytes, Index + IndexShift); - break; - } - case kernel_param_kind_t::kind_work_group_memory: { - addArg(kernel_param_kind_t::kind_std_layout, nullptr, - static_cast(Ptr)->buffer_size, - Index + IndexShift); - break; - } - case kernel_param_kind_t::kind_sampler: { - addArg(kernel_param_kind_t::kind_sampler, Ptr, sizeof(sampler), - Index + IndexShift); - break; - } - case kernel_param_kind_t::kind_specialization_constants_buffer: { - addArg(kernel_param_kind_t::kind_specialization_constants_buffer, Ptr, Size, - Index + IndexShift); - break; - } - case kernel_param_kind_t::kind_invalid: - throw exception(make_error_code(errc::invalid), - "Invalid kernel param kind"); - break; - } + impl->MKernelData.processArg(Ptr, Kind, Size, Index, IndexShift, + IsKernelCreatedFromSource, impl->MNDRDesc); } +#endif void handler::setArgHelper(int ArgIndex, detail::work_group_memory_impl &Arg) { impl->MWorkGroupMemoryObjects.push_back( @@ -1309,102 +1076,26 @@ void handler::setArgHelper(int ArgIndex, stream &&Str) { ArgIndex); } -// The argument can take up more space to store additional information about -// MAccessRange, MMemoryRange, and MOffset added with addArgsForGlobalAccessor. -// We use the worst-case estimate because the lifetime of the vector is short. -// In processArg the kind_stream case introduces the maximum number of -// additional arguments. The case adds additional 12 arguments to the currently -// processed argument, hence worst-case estimate is 12+1=13. -// TODO: the constant can be removed if the size of MArgs will be calculated at -// compile time. -inline constexpr size_t MaxNumAdditionalArgs = 13; - void handler::extractArgsAndReqs() { assert(MKernel && "MKernel is not initialized"); - std::vector UnPreparedArgs = std::move(impl->MArgs); - clearArgs(); - - std::sort( - UnPreparedArgs.begin(), UnPreparedArgs.end(), - [](const detail::ArgDesc &first, const detail::ArgDesc &second) -> bool { - return (first.MIndex < second.MIndex); - }); - - const bool IsKernelCreatedFromSource = MKernel->isCreatedFromSource(); - impl->MArgs.reserve(MaxNumAdditionalArgs * UnPreparedArgs.size()); - - size_t IndexShift = 0; - for (size_t I = 0; I < UnPreparedArgs.size(); ++I) { - void *Ptr = UnPreparedArgs[I].MPtr; - const detail::kernel_param_kind_t &Kind = UnPreparedArgs[I].MType; - const int &Size = UnPreparedArgs[I].MSize; - const int Index = UnPreparedArgs[I].MIndex; - processArg(Ptr, Kind, Size, Index, IndexShift, IsKernelCreatedFromSource, - false); - } + impl->MKernelData.extractArgsAndReqs(impl->MNDRDesc, + MKernel->isCreatedFromSource()); } +#ifndef __INTEL_PREVIEW_BREAKING_CHANGES +// TODO: Those functions are not used anymore, remove it in the next +// ABI-breaking window. void handler::extractArgsAndReqsFromLambda( char *LambdaPtr, detail::kernel_param_desc_t (*ParamDescGetter)(int), size_t NumKernelParams, bool IsESIMD) { - size_t IndexShift = 0; - impl->MArgs.reserve(MaxNumAdditionalArgs * NumKernelParams); - - for (size_t I = 0; I < NumKernelParams; ++I) { - detail::kernel_param_desc_t ParamDesc = ParamDescGetter(I); - void *Ptr = LambdaPtr + ParamDesc.offset; - const detail::kernel_param_kind_t &Kind = ParamDesc.kind; - const int &Size = ParamDesc.info; - if (Kind == detail::kernel_param_kind_t::kind_accessor) { - // For args kind of accessor Size is information about accessor. - // The first 11 bits of Size encodes the accessor target. - const access::target AccTarget = - static_cast(Size & AccessTargetMask); - if ((AccTarget == access::target::device || - AccTarget == access::target::constant_buffer) || - (AccTarget == access::target::image || - AccTarget == access::target::image_array)) { - detail::AccessorBaseHost *AccBase = - static_cast(Ptr); - Ptr = detail::getSyclObjImpl(*AccBase).get(); - } else if (AccTarget == access::target::local) { - detail::LocalAccessorBaseHost *LocalAccBase = - static_cast(Ptr); - Ptr = detail::getSyclObjImpl(*LocalAccBase).get(); - } - } else if (Kind == detail::kernel_param_kind_t::kind_dynamic_accessor) { - // For args kind of accessor Size is information about accessor. - // The first 11 bits of Size encodes the accessor target. - // Only local targets are supported for dynamic accessors. - assert(static_cast(Size & AccessTargetMask) == - access::target::local); - - ext::oneapi::experimental::detail::dynamic_parameter_base - *DynamicParamBase = static_cast< - ext::oneapi::experimental::detail::dynamic_parameter_base *>(Ptr); - Ptr = detail::getSyclObjImpl(*DynamicParamBase).get(); - } else if (Kind == - detail::kernel_param_kind_t::kind_dynamic_work_group_memory) { - ext::oneapi::experimental::detail::dynamic_parameter_base - *DynamicParamBase = static_cast< - ext::oneapi::experimental::detail::dynamic_parameter_base *>(Ptr); - Ptr = detail::getSyclObjImpl(*DynamicParamBase).get(); - } - - processArg(Ptr, Kind, Size, I, IndexShift, - /*IsKernelCreatedFromSource=*/false, IsESIMD); - } + impl->MKernelData.extractArgsAndReqsFromLambda(impl->MNDRDesc); } -#ifndef __INTEL_PREVIEW_BREAKING_CHANGES -// TODO: Those functions are not used anymore, remove it in the next -// ABI-breaking window. void handler::extractArgsAndReqsFromLambda( char *LambdaPtr, const std::vector &ParamDescs, bool IsESIMD) { const bool IsKernelCreatedFromSource = false; size_t IndexShift = 0; - impl->MArgs.reserve(MaxNumAdditionalArgs * ParamDescs.size()); for (size_t I = 0; I < ParamDescs.size(); ++I) { void *Ptr = LambdaPtr + ParamDescs[I].offset; @@ -1428,8 +1119,8 @@ void handler::extractArgsAndReqsFromLambda( Ptr = detail::getSyclObjImpl(*LocalAccBase).get(); } } - processArg(Ptr, Kind, Size, I, IndexShift, IsKernelCreatedFromSource, - IsESIMD); + impl->MKernelData.processArg(Ptr, Kind, Size, I, IndexShift, + IsKernelCreatedFromSource, impl->MNDRDesc); } } @@ -2447,7 +2138,7 @@ void handler::registerDynamicParameter( "Dynamic Parameters cannot be used with normal SYCL submissions"); } - impl->MDynamicParameters.emplace_back(DynamicParamImpl, ArgIndex); + impl->MKernelData.addDynamicParameter(DynamicParamImpl, ArgIndex); } #ifndef __INTEL_PREVIEW_BREAKING_CHANGES @@ -2500,13 +2191,15 @@ void handler::addLifetimeSharedPtrStorage(std::shared_ptr SPtr) { void handler::addArg(detail::kernel_param_kind_t ArgKind, void *Req, int AccessTarget, int ArgIndex) { - impl->MArgs.emplace_back(ArgKind, Req, AccessTarget, ArgIndex); + impl->MKernelData.addArg(ArgKind, Req, AccessTarget, ArgIndex); } -void handler::clearArgs() { impl->MArgs.clear(); } +#ifndef __INTEL_PREVIEW_BREAKING_CHANGES +void handler::clearArgs() { impl->MKernelData.clearArgs(); } +#endif void handler::setArgsToAssociatedAccessors() { - impl->MArgs = impl->MAssociatedAccesors; + impl->MKernelData.setArgs(impl->MAssociatedAccesors); } bool handler::HasAssociatedAccessor(detail::AccessorImplHost *Req, diff --git a/sycl/unittests/program_manager/arg_mask/EliminatedArgMask.cpp b/sycl/unittests/program_manager/arg_mask/EliminatedArgMask.cpp index aea61760c1e70..4153eabffeaa3 100644 --- a/sycl/unittests/program_manager/arg_mask/EliminatedArgMask.cpp +++ b/sycl/unittests/program_manager/arg_mask/EliminatedArgMask.cpp @@ -147,7 +147,7 @@ class MockHandler : public sycl::handler { CommandGroup.reset(new sycl::detail::CGExecKernel( std::move(impl->MNDRDesc), std::move(CGH->MHostKernel), std::move(CGH->MKernel), std::move(impl->MKernelBundle), - std::move(impl->CGData), std::move(impl->MArgs), + std::move(impl->CGData), std::move(impl->MKernelData).getArgs(), CGH->MKernelName.data(), *impl->MKernelData.getDeviceKernelInfo(), std::move(CGH->MStreamStorage), std::move(impl->MAuxiliaryResources), impl->MCGType, {}, impl->MKernelIsCooperative, diff --git a/sycl/unittests/scheduler/SchedulerTestUtils.hpp b/sycl/unittests/scheduler/SchedulerTestUtils.hpp index 067d3d201e2dd..fec0f88327e30 100644 --- a/sycl/unittests/scheduler/SchedulerTestUtils.hpp +++ b/sycl/unittests/scheduler/SchedulerTestUtils.hpp @@ -252,7 +252,9 @@ class MockHandler : public sycl::handler { std::vector &getEvents() { return impl->CGData.MEvents; } - std::vector &getArgs() { return impl->MArgs; } + std::vector &getArgs() { + return impl->MKernelData.getArgs(); + } sycl::detail::KernelNameStrT getKernelName() { return toKernelNameStrT(MKernelName); } From 971672b314ab71cacf92b694b4dff6ddb9c0b918 Mon Sep 17 00:00:00 2001 From: Sergei Vinogradov Date: Wed, 20 Aug 2025 23:07:57 +0200 Subject: [PATCH 3/6] Fix unused parameters --- sycl/source/handler.cpp | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index 2986b30120da8..35e63f334e948 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -1058,6 +1058,7 @@ void handler::associateWithHandler( void handler::processArg(void *Ptr, const detail::kernel_param_kind_t &Kind, const int Size, const size_t Index, size_t &IndexShift, bool IsKernelCreatedFromSource, bool IsESIMD) { + (void)IsESIMD; impl->MKernelData.processArg(Ptr, Kind, Size, Index, IndexShift, IsKernelCreatedFromSource, impl->MNDRDesc); } @@ -1088,12 +1089,17 @@ void handler::extractArgsAndReqs() { void handler::extractArgsAndReqsFromLambda( char *LambdaPtr, detail::kernel_param_desc_t (*ParamDescGetter)(int), size_t NumKernelParams, bool IsESIMD) { + (void)LambdaPtr; + (void)ParamDescGetter; + (void)NumKernelParams; + (void)IsESIMD; impl->MKernelData.extractArgsAndReqsFromLambda(impl->MNDRDesc); } void handler::extractArgsAndReqsFromLambda( char *LambdaPtr, const std::vector &ParamDescs, bool IsESIMD) { + (void)IsESIMD; const bool IsKernelCreatedFromSource = false; size_t IndexShift = 0; From a1c85b24b44beabb22c66e88464404f40280da29 Mon Sep 17 00:00:00 2001 From: Sergei Vinogradov Date: Thu, 21 Aug 2025 17:09:40 +0200 Subject: [PATCH 4/6] Fix backward compatibility issue found by CI --- sycl/source/detail/kernel_data.hpp | 6 ++++++ sycl/source/handler.cpp | 2 +- 2 files changed, 7 insertions(+), 1 deletion(-) diff --git a/sycl/source/detail/kernel_data.hpp b/sycl/source/detail/kernel_data.hpp index f1891347cc9ab..098d87a167355 100644 --- a/sycl/source/detail/kernel_data.hpp +++ b/sycl/source/detail/kernel_data.hpp @@ -69,6 +69,12 @@ class KernelData { return MDeviceKernelInfoPtr->ParamDescGetter; } +#ifndef __INTEL_PREVIEW_BREAKING_CHANGES + // TODO: remove this method in the next ABI-breaking window + // it is used by handler code that will be removed in the next + // ABI-breaking window + void setESIMD(bool IsESIMD) { MDeviceKernelInfoPtr->IsESIMD = IsESIMD; } +#endif bool isESIMD() const { return MDeviceKernelInfoPtr->IsESIMD; } bool hasSpecialCaptures() const { diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index 35e63f334e948..9eaf56529735c 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -1099,9 +1099,9 @@ void handler::extractArgsAndReqsFromLambda( void handler::extractArgsAndReqsFromLambda( char *LambdaPtr, const std::vector &ParamDescs, bool IsESIMD) { - (void)IsESIMD; const bool IsKernelCreatedFromSource = false; size_t IndexShift = 0; + impl->MKernelData.setESIMD(IsESIMD); for (size_t I = 0; I < ParamDescs.size(); ++I) { void *Ptr = LambdaPtr + ParamDescs[I].offset; From d6cc396a8740a3acf11085f85b55c6273ab48ba9 Mon Sep 17 00:00:00 2001 From: Sergei Vinogradov Date: Thu, 28 Aug 2025 18:45:53 +0200 Subject: [PATCH 5/6] Fix issue after rebase --- sycl/source/detail/kernel_data.hpp | 1 - 1 file changed, 1 deletion(-) diff --git a/sycl/source/detail/kernel_data.hpp b/sycl/source/detail/kernel_data.hpp index 098d87a167355..b8ca959733b52 100644 --- a/sycl/source/detail/kernel_data.hpp +++ b/sycl/source/detail/kernel_data.hpp @@ -13,7 +13,6 @@ #include #include -#include #include From 6163e97cd0c7e3c485ea4225999e5eb943b5f186 Mon Sep 17 00:00:00 2001 From: Sergei Vinogradov Date: Thu, 28 Aug 2025 22:55:10 +0200 Subject: [PATCH 6/6] Fix unit tests after rebase --- sycl/include/sycl/handler.hpp | 2 -- sycl/test/abi/sycl_symbols_linux.dump | 1 + sycl/unittests/program_manager/arg_mask/EliminatedArgMask.cpp | 2 +- 3 files changed, 2 insertions(+), 3 deletions(-) diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 3d81825ab831b..ab36c943ed128 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -851,8 +851,6 @@ class __SYCL_EXPORT handler { if constexpr (KernelHasName) { // TODO support ESIMD in no-integration-header case too. - // Force hasSpecialCaptures to be evaluated at compile-time. - constexpr bool HasSpecialCapt = detail::hasSpecialCaptures(); setKernelInfo((void *)MHostKernel->getPtr(), &detail::getDeviceKernelInfo()); diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index fd08eab77fdf3..746a08b47adbb 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3546,6 +3546,7 @@ _ZN4sycl3_V17handler12setArgHelperEiONS0_6streamE _ZN4sycl3_V17handler12setArgHelperEiRNS0_6detail22work_group_memory_implE _ZN4sycl3_V17handler13getKernelNameEv _ZN4sycl3_V17handler13setKernelInfoEPviPFNS0_6detail19kernel_param_desc_tEiEbb +_ZN4sycl3_V17handler13setKernelInfoEPvPNS0_6detail16DeviceKernelInfoE _ZN4sycl3_V17handler14addAccessorReqESt10shared_ptrINS0_6detail16AccessorImplHostEE _ZN4sycl3_V17handler14setNDRangeUsedEb _ZN4sycl3_V17handler15ext_oneapi_copyENS0_3ext6oneapi12experimental16image_mem_handleENS0_5rangeILi3EEERKNS4_16image_descriptorEPvS7_S7_S7_ diff --git a/sycl/unittests/program_manager/arg_mask/EliminatedArgMask.cpp b/sycl/unittests/program_manager/arg_mask/EliminatedArgMask.cpp index 4153eabffeaa3..66c06794992c9 100644 --- a/sycl/unittests/program_manager/arg_mask/EliminatedArgMask.cpp +++ b/sycl/unittests/program_manager/arg_mask/EliminatedArgMask.cpp @@ -148,7 +148,7 @@ class MockHandler : public sycl::handler { std::move(impl->MNDRDesc), std::move(CGH->MHostKernel), std::move(CGH->MKernel), std::move(impl->MKernelBundle), std::move(impl->CGData), std::move(impl->MKernelData).getArgs(), - CGH->MKernelName.data(), *impl->MKernelData.getDeviceKernelInfo(), + CGH->MKernelName.data(), *impl->MKernelData.getDeviceKernelInfoPtr(), std::move(CGH->MStreamStorage), std::move(impl->MAuxiliaryResources), impl->MCGType, {}, impl->MKernelIsCooperative, impl->MKernelUsesClusterLaunch, impl->MKernelWorkGroupMemorySize,