diff --git a/sycl/doc/extensions/LevelZeroBackend/LevelZeroBackend.md b/sycl/doc/extensions/LevelZeroBackend/LevelZeroBackend.md index 55ccf83219f1e..ca6b1606b2f89 100644 --- a/sycl/doc/extensions/LevelZeroBackend/LevelZeroBackend.md +++ b/sycl/doc/extensions/LevelZeroBackend/LevelZeroBackend.md @@ -90,7 +90,8 @@ a SYCL object that encapsulates a corresponding Level-Zero object: |``` make(const vector_class &, ze_context_handle_t, ownership = transfer);```| Constructs a SYCL context instance from a Level-Zero ```ze_context_handle_t```. The context is created against the devices passed in. There must be at least one device given and all the devices must be from the same SYCL platform and thus from the same Level-Zero driver. The ```ownership``` argument specifies if the SYCL runtime should take ownership of the passed native handle. The default behavior is to transfer the ownership to the SYCL runtime. See section 4.4 for details.| |``` make(const context &, ze_command_queue_handle_t, ownership = transfer);```| Constructs a SYCL queue instance from a Level-Zero ```ze_command_queue_handle_t```. The context argument must be a valid SYCL context encapsulating a Level-Zero context. The queue is attached to the first device in the passed SYCL context. The ```ownership``` argument specifies if the SYCL runtime should take ownership of the passed native handle. The default behavior is to transfer the ownership to the SYCL runtime. See section 4.4 for details.| |``` make(const context &, ze_event_handle_t, ownership = transfer);```| Constructs a SYCL event instance from a Level-Zero ```ze_event_handle_t```. The context argument must be a valid SYCL context encapsulating a Level-Zero context. The Level-Zero event should be allocated from an event pool created in the same context. The ```ownership``` argument specifies if the SYCL runtime should take ownership of the passed native handle. The default behavior is to transfer the ownership to the SYCL runtime. See section 4.4 for details.| -|``` make(const context &, ze_module_handle_t);```| Constructs a SYCL program instance from a Level-Zero ```ze_module_handle_t```. The context argument must be a valid SYCL context encapsulating a Level-Zero context. The Level-Zero module must be fully linked (i.e. not require further linking through [```zeModuleDynamicLink```](https://spec.oneapi.com/level-zero/latest/core/api.html?highlight=zemoduledynamiclink#_CPPv419zeModuleDynamicLink8uint32_tP18ze_module_handle_tP28ze_module_build_log_handle_t)), and thus the SYCL program is created in the "linked" state.| +|``` make(const context &, ze_kernel_handle_t, ownership = transfer);```| Constructs a SYCL kernel from a Level-Zero ```ze_kernel_handle_t```. The context argument must be a valid SYCL context encapsulating a Level-Zero context. The ```ownership``` argument specifies if the SYCL runtime should take ownership of the passed native handle. The default behavior is to transfer the ownership to the SYCL runtime. See section 4.4 for details.| +|``` make(const context &, ze_module_handle_t, ownership = transfer);```| Constructs a SYCL program instance from a Level-Zero ```ze_module_handle_t```. The context argument must be a valid SYCL context encapsulating a Level-Zero context. The Level-Zero module must be fully linked (i.e. not require further linking through [```zeModuleDynamicLink```](https://spec.oneapi.com/level-zero/latest/core/api.html?highlight=zemoduledynamiclink#_CPPv419zeModuleDynamicLink8uint32_tP18ze_module_handle_tP28ze_module_build_log_handle_t)), and thus the SYCL program is created in the "linked" state. The ```ownership``` argument specifies if the SYCL runtime should take ownership of the passed native handle. The default behavior is to transfer the ownership to the SYCL runtime. See section 4.4 for details.| NOTE: We shall consider adding other interoperability as needed, if possible. @@ -197,3 +198,5 @@ struct free_memory { |4|2021-07-06|Rehana Begam|Introduced explicit ownership for queue |5|2021-07-25|Sergey Maslov|Introduced SYCL interop for events |6|2021-08-30|Dmitry Vodopyanov|Updated according to SYCL 2020 reqs for extensions +|7|2021-09-09|Rehana Begam|Introduced explicit ownership for kernel_bundle +|8|2021-09-10|Rehana Begam|Introduced explicit ownership for kernel diff --git a/sycl/include/CL/sycl/backend.hpp b/sycl/include/CL/sycl/backend.hpp index c4a21266ec1f4..5a795b9e1b1aa 100644 --- a/sycl/include/CL/sycl/backend.hpp +++ b/sycl/include/CL/sycl/backend.hpp @@ -108,6 +108,12 @@ __SYCL_EXPORT event make_event(pi_native_handle NativeHandle, backend Backend); __SYCL_EXPORT kernel make_kernel(pi_native_handle NativeHandle, const context &TargetContext, backend Backend); +__SYCL_EXPORT kernel make_kernel(pi_native_handle NativeHandle, + const context &TargetContext, + bool KeepOwnership, backend Backend); +__SYCL_EXPORT std::shared_ptr +make_kernel_bundle(pi_native_handle NativeHandle, const context &TargetContext, + bool KeepOwnership, bundle_state State, backend Backend); __SYCL_EXPORT std::shared_ptr make_kernel_bundle(pi_native_handle NativeHandle, const context &TargetContext, bundle_state State, backend Backend); @@ -199,13 +205,36 @@ make_buffer(const typename backend_traits::template input_type< reinterpret_cast(BackendObject), TargetContext, AvailableEvent); } +template +kernel +make_kernel(const typename backend_traits::template input_type + &BackendObject, + const context &TargetContext, bool KeepOwnership) { + return detail::make_kernel(detail::pi::cast(BackendObject), + TargetContext, KeepOwnership, Backend); +} + template kernel make_kernel(const typename backend_traits::template input_type &BackendObject, const context &TargetContext) { return detail::make_kernel(detail::pi::cast(BackendObject), - TargetContext, Backend); + TargetContext, false, Backend); +} + +template +typename std::enable_if< + detail::InteropFeatureSupportMap::MakeKernelBundle == true, + kernel_bundle>::type +make_kernel_bundle(const typename backend_traits::template input_type< + kernel_bundle> &BackendObject, + const context &TargetContext, bool KeepOwnership) { + std::shared_ptr KBImpl = + detail::make_kernel_bundle( + detail::pi::cast(BackendObject), TargetContext, + KeepOwnership, State, Backend); + return detail::createSyclObjFromImpl>(KBImpl); } template @@ -218,7 +247,7 @@ make_kernel_bundle(const typename backend_traits::template input_type< std::shared_ptr KBImpl = detail::make_kernel_bundle( detail::pi::cast(BackendObject), TargetContext, - State, Backend); + false, State, Backend); return detail::createSyclObjFromImpl>(KBImpl); } } // namespace sycl diff --git a/sycl/include/CL/sycl/detail/pi.h b/sycl/include/CL/sycl/detail/pi.h index 3ab941823e603..24ccae9847a14 100644 --- a/sycl/include/CL/sycl/detail/pi.h +++ b/sycl/include/CL/sycl/detail/pi.h @@ -40,10 +40,14 @@ // changes the API version from 3.5 to 4.6. // 5.7 Added new context and ownership arguments to // piextEventCreateWithNativeHandle +// 6.8 Added new ownership argument to piextProgramCreateWithNativeHandle which +// changes the API version from 5.7 to 6.8 +// 7.9 Added new ownership argument to piextKernelCreateWithNativeHandle which +// changes the API version from 6.8 to 7.9 // #include "CL/cl.h" -#define _PI_H_VERSION_MAJOR 5 -#define _PI_H_VERSION_MINOR 7 +#define _PI_H_VERSION_MAJOR 7 +#define _PI_H_VERSION_MINOR 9 #define _PI_STRING_HELPER(a) #a #define _PI_CONCAT(a, b) _PI_STRING_HELPER(a.b) @@ -1220,8 +1224,11 @@ piextProgramGetNativeHandle(pi_program program, pi_native_handle *nativeHandle); /// \param nativeHandle is the native handle to create PI program from. /// \param context is the PI context of the program. /// \param program is the PI program created from the native handle. +/// \param ownNativeHandle tells if SYCL RT should assume the ownership of +/// the native handle, if it can. __SYCL_EXPORT pi_result piextProgramCreateWithNativeHandle( - pi_native_handle nativeHandle, pi_context context, pi_program *program); + pi_native_handle nativeHandle, pi_context context, pi_program *program, + bool ownNativeHandle); // // Kernel @@ -1315,12 +1322,13 @@ __SYCL_EXPORT pi_result piKernelSetExecInfo(pi_kernel kernel, /// /// \param nativeHandle is the native handle to create PI kernel from. /// \param context is the PI context of the kernel. +/// \param kernel is the PI kernel created from the native handle. /// \param ownNativeHandle tells if SYCL RT should assume the ownership of /// the native handle, if it can. -/// \param kernel is the PI kernel created from the native handle. +/// __SYCL_EXPORT pi_result piextKernelCreateWithNativeHandle( - pi_native_handle nativeHandle, pi_context context, bool ownNativeHandle, - pi_kernel *kernel); + pi_native_handle nativeHandle, pi_context context, pi_kernel *kernel, + bool ownNativeHandle); /// Gets the native handle of a PI kernel object. /// diff --git a/sycl/plugins/cuda/pi_cuda.cpp b/sycl/plugins/cuda/pi_cuda.cpp index aec05da87982d..c475899ecb43c 100644 --- a/sycl/plugins/cuda/pi_cuda.cpp +++ b/sycl/plugins/cuda/pi_cuda.cpp @@ -2681,7 +2681,7 @@ pi_result cuda_piEnqueueNativeKernel(pi_queue, void (*)(void *), void *, size_t, } pi_result cuda_piextKernelCreateWithNativeHandle(pi_native_handle, pi_context, - bool, pi_kernel *) { + pi_kernel *, bool) { sycl::detail::pi::die("Unsupported operation"); return PI_SUCCESS; } @@ -3161,10 +3161,14 @@ pi_result cuda_piextProgramGetNativeHandle(pi_program program, /// \param[in] nativeHandle The native handle to create PI program object from. /// \param[in] context The PI context of the program. /// \param[out] program Set to the PI program object created from native handle. +/// \param[in] ownNativeHandle tells if SYCL RT should assume the ownership of +/// the native handle, if it can. /// /// \return TBD pi_result cuda_piextProgramCreateWithNativeHandle(pi_native_handle, pi_context, - pi_program *) { + pi_program *, + bool ownNativeHandle) { + (void)ownNativeHandle; cl::sycl::detail::pi::die( "Creation of PI program from native handle not implemented"); return {}; diff --git a/sycl/plugins/esimd_cpu/pi_esimd_cpu.cpp b/sycl/plugins/esimd_cpu/pi_esimd_cpu.cpp index 1035c05e23caa..0ecc5933c3e30 100644 --- a/sycl/plugins/esimd_cpu/pi_esimd_cpu.cpp +++ b/sycl/plugins/esimd_cpu/pi_esimd_cpu.cpp @@ -1027,7 +1027,7 @@ pi_result piextProgramGetNativeHandle(pi_program, pi_native_handle *) { } pi_result piextProgramCreateWithNativeHandle(pi_native_handle, pi_context, - pi_program *) { + pi_program *, bool) { DIE_NO_IMPLEMENTATION; } @@ -1397,8 +1397,8 @@ piEnqueueKernelLaunch(pi_queue Queue, pi_kernel Kernel, pi_uint32 WorkDim, return PI_SUCCESS; } -pi_result piextKernelCreateWithNativeHandle(pi_native_handle, pi_context, bool, - pi_kernel *) { +pi_result piextKernelCreateWithNativeHandle(pi_native_handle, pi_context, + pi_kernel *, bool) { DIE_NO_IMPLEMENTATION; } diff --git a/sycl/plugins/level_zero/pi_level_zero.cpp b/sycl/plugins/level_zero/pi_level_zero.cpp index 338449afcda52..b864bac0464ad 100644 --- a/sycl/plugins/level_zero/pi_level_zero.cpp +++ b/sycl/plugins/level_zero/pi_level_zero.cpp @@ -3235,7 +3235,7 @@ pi_result piProgramCreate(pi_context Context, const void *ILBytes, // and piProgramCompile. Also it is only then we know the build options. try { - *Program = new _pi_program(Context, ILBytes, Length, _pi_program::IL); + *Program = new _pi_program(Context, ILBytes, Length, _pi_program::IL, true); } catch (const std::bad_alloc &) { return PI_OUT_OF_HOST_MEMORY; } catch (...) { @@ -3281,7 +3281,8 @@ pi_result piProgramCreateWithBinary( // information to distinguish the cases. try { - *Program = new _pi_program(Context, Binary, Length, _pi_program::Native); + *Program = + new _pi_program(Context, Binary, Length, _pi_program::Native, true); } catch (const std::bad_alloc &) { return PI_OUT_OF_HOST_MEMORY; } catch (...) { @@ -3528,7 +3529,7 @@ pi_result piProgramLink(pi_context Context, pi_uint32 NumDevices, return res; } Input = new _pi_program(Input->Context, ZeModule, _pi_program::Object, - Input->HasImports); + true, Input->HasImports); Input->HasImportsAndIsLinked = true; } } else { @@ -3551,7 +3552,8 @@ pi_result piProgramLink(pi_context Context, pi_uint32 NumDevices, // the description of the failure). if (ZeResult == ZE_RESULT_SUCCESS || ZeResult == ZE_RESULT_ERROR_MODULE_LINK_FAILURE) { - *RetProgram = new _pi_program(Context, std::move(Inputs), ZeBuildLog); + *RetProgram = + new _pi_program(Context, std::move(Inputs), ZeBuildLog, true); } if (ZeResult != ZE_RESULT_SUCCESS) return mapError(ZeResult); @@ -3783,7 +3785,8 @@ pi_result piextProgramGetNativeHandle(pi_program Program, pi_result piextProgramCreateWithNativeHandle(pi_native_handle NativeHandle, pi_context Context, - pi_program *Program) { + pi_program *Program, + bool OwnNativeHandle) { PI_ASSERT(Program, PI_INVALID_PROGRAM); PI_ASSERT(NativeHandle, PI_INVALID_VALUE); PI_ASSERT(Context, PI_INVALID_CONTEXT); @@ -3795,7 +3798,8 @@ pi_result piextProgramCreateWithNativeHandle(pi_native_handle NativeHandle, // executable (state Object). try { - *Program = new _pi_program(Context, ZeModule, _pi_program::Exe); + *Program = + new _pi_program(Context, ZeModule, _pi_program::Exe, OwnNativeHandle); } catch (const std::bad_alloc &) { return PI_OUT_OF_HOST_MEMORY; } catch (...) { @@ -4352,8 +4356,8 @@ piEnqueueKernelLaunch(pi_queue Queue, pi_kernel Kernel, pi_uint32 WorkDim, return PI_SUCCESS; } -pi_result piextKernelCreateWithNativeHandle(pi_native_handle, pi_context, bool, - pi_kernel *) { +pi_result piextKernelCreateWithNativeHandle(pi_native_handle, pi_context, + pi_kernel *, bool) { die("Unsupported operation"); return PI_SUCCESS; } diff --git a/sycl/plugins/level_zero/pi_level_zero.hpp b/sycl/plugins/level_zero/pi_level_zero.hpp index 748fb331d9025..1ab112fd320d9 100644 --- a/sycl/plugins/level_zero/pi_level_zero.hpp +++ b/sycl/plugins/level_zero/pi_level_zero.hpp @@ -1056,25 +1056,27 @@ struct _pi_program : _pi_object { }; // Construct a program in IL or Native state. - _pi_program(pi_context Context, const void *Input, size_t Length, state St) - : State(St), Context(Context), Code(new uint8_t[Length]), - CodeLength(Length), ZeModule(nullptr), HasImports(false), - HasImportsAndIsLinked(false), ZeBuildLog(nullptr) { + _pi_program(pi_context Context, const void *Input, size_t Length, state St, + bool OwnZeModule) + : State(St), Context(Context), OwnZeModule(OwnZeModule), + Code(new uint8_t[Length]), CodeLength(Length), ZeModule(nullptr), + HasImports(false), HasImportsAndIsLinked(false), ZeBuildLog(nullptr) { std::memcpy(Code.get(), Input, Length); } // Construct a program in either Object or Exe state. _pi_program(pi_context Context, ze_module_handle_t ZeModule, state St, - bool HasImports = false) - : State(St), Context(Context), ZeModule(ZeModule), HasImports(HasImports), + bool OwnZeModule, bool HasImports = false) + : State(St), Context(Context), OwnZeModule(OwnZeModule), + ZeModule(ZeModule), HasImports(HasImports), HasImportsAndIsLinked(false), ZeBuildLog(nullptr) {} // Construct a program in LinkedExe state. _pi_program(pi_context Context, std::vector &&Inputs, - ze_module_build_log_handle_t ZeLog) - : State(LinkedExe), Context(Context), ZeModule(nullptr), - HasImports(false), HasImportsAndIsLinked(false), + ze_module_build_log_handle_t ZeLog, bool OwnZeModule) + : State(LinkedExe), Context(Context), OwnZeModule(OwnZeModule), + ZeModule(nullptr), HasImports(false), HasImportsAndIsLinked(false), LinkedPrograms(std::move(Inputs)), ZeBuildLog(ZeLog) {} ~_pi_program(); @@ -1083,6 +1085,10 @@ struct _pi_program : _pi_object { state State; pi_context Context; // Context of the program. + // Indicates if we own the ZeModule or it came from interop that + // asked to not transfer the ownership to SYCL RT. + bool OwnZeModule; + // Used for programs in IL or Native states. std::unique_ptr Code; // Array containing raw IL / native code. size_t CodeLength; // Size (bytes) of the array. diff --git a/sycl/plugins/opencl/pi_opencl.cpp b/sycl/plugins/opencl/pi_opencl.cpp index fb05ceca94b32..7ebbaaf86d1e5 100644 --- a/sycl/plugins/opencl/pi_opencl.cpp +++ b/sycl/plugins/opencl/pi_opencl.cpp @@ -444,8 +444,9 @@ pi_result piProgramCreate(pi_context context, const void *il, size_t length, } pi_result piextProgramCreateWithNativeHandle(pi_native_handle nativeHandle, - pi_context, - pi_program *piProgram) { + pi_context, pi_program *piProgram, + bool ownNativeHandle) { + (void)ownNativeHandle; assert(piProgram != nullptr); *piProgram = reinterpret_cast(nativeHandle); return PI_SUCCESS; @@ -497,8 +498,9 @@ pi_result piextKernelSetArgSampler(pi_kernel kernel, pi_uint32 arg_index, } pi_result piextKernelCreateWithNativeHandle(pi_native_handle nativeHandle, - pi_context, bool, - pi_kernel *piKernel) { + pi_context, pi_kernel *piKernel, + bool ownNativeHandle) { + (void)ownNativeHandle; assert(piKernel != nullptr); *piKernel = reinterpret_cast(nativeHandle); return PI_SUCCESS; diff --git a/sycl/source/backend.cpp b/sycl/source/backend.cpp index ef68f694ec4a9..e8640c6fc859c 100644 --- a/sycl/source/backend.cpp +++ b/sycl/source/backend.cpp @@ -119,12 +119,18 @@ __SYCL_EXPORT event make_event(pi_native_handle NativeHandle, std::shared_ptr make_kernel_bundle(pi_native_handle NativeHandle, const context &TargetContext, bundle_state State, backend Backend) { + return make_kernel_bundle(NativeHandle, TargetContext, false, State, Backend); +} + +std::shared_ptr +make_kernel_bundle(pi_native_handle NativeHandle, const context &TargetContext, + bool KeepOwnership, bundle_state State, backend Backend) { const auto &Plugin = getPlugin(Backend); const auto &ContextImpl = getSyclObjImpl(TargetContext); pi::PiProgram PiProgram = nullptr; Plugin.call( - NativeHandle, ContextImpl->getHandleRef(), &PiProgram); + NativeHandle, ContextImpl->getHandleRef(), &PiProgram, !KeepOwnership); std::vector ProgramDevices; size_t NumDevices = 0; @@ -194,14 +200,20 @@ make_kernel_bundle(pi_native_handle NativeHandle, const context &TargetContext, return std::make_shared(TargetContext, Devices, DevImg); } + kernel make_kernel(pi_native_handle NativeHandle, const context &TargetContext, backend Backend) { + return make_kernel(NativeHandle, TargetContext, false, Backend); +} + +kernel make_kernel(pi_native_handle NativeHandle, const context &TargetContext, + bool KeepOwnership, backend Backend) { const auto &Plugin = getPlugin(Backend); const auto &ContextImpl = getSyclObjImpl(TargetContext); // Create PI kernel first. pi::PiKernel PiKernel = nullptr; Plugin.call( - NativeHandle, ContextImpl->getHandleRef(), false, &PiKernel); + NativeHandle, ContextImpl->getHandleRef(), &PiKernel, !KeepOwnership); if (Backend == backend::opencl) Plugin.call(PiKernel); diff --git a/sycl/source/backend/level_zero.cpp b/sycl/source/backend/level_zero.cpp index fe10236415ec0..f65eac9a4eedc 100644 --- a/sycl/source/backend/level_zero.cpp +++ b/sycl/source/backend/level_zero.cpp @@ -95,6 +95,22 @@ __SYCL_EXPORT queue make_queue(const context &Context, return make_queue(Context, NativeHandle, false); } +//---------------------------------------------------------------------------- +// Implementation of level_zero::make +__SYCL_EXPORT kernel make_kernel(const context &Context, + pi_native_handle NativeHandle, + bool KeepOwnership) { + const auto &ContextImpl = getSyclObjImpl(Context); + return detail::make_kernel(NativeHandle, Context, KeepOwnership, + backend::level_zero); +} + +// TODO: remove this version (without ownership) when allowed to break ABI. +__SYCL_EXPORT kernel make_kernel(const context &Context, + pi_native_handle NativeHandle) { + return make_kernel(Context, NativeHandle, false); +} + //---------------------------------------------------------------------------- // Implementation of level_zero::make __SYCL_EXPORT event make_event(const context &Context, diff --git a/sycl/source/backend/opencl.cpp b/sycl/source/backend/opencl.cpp index 7e76478c8e630..7a6d29791d0e2 100644 --- a/sycl/source/backend/opencl.cpp +++ b/sycl/source/backend/opencl.cpp @@ -55,6 +55,14 @@ __SYCL_EXPORT queue make_queue(const context &Context, return detail::make_queue(NativeHandle, Context, false, ContextImpl->get_async_handler(), backend::opencl); } + +//---------------------------------------------------------------------------- +// Implementation of opencl::make +__SYCL_EXPORT kernel make_kernel(const context &Context, + pi_native_handle NativeHandle) { + const auto &ContextImpl = getSyclObjImpl(Context); + return detail::make_kernel(NativeHandle, Context, false, backend::opencl); +} } // namespace opencl } // namespace sycl } // __SYCL_INLINE_NAMESPACE(cl) diff --git a/sycl/source/detail/program_impl.cpp b/sycl/source/detail/program_impl.cpp index 8090280afa7e0..2a788ae62514c 100644 --- a/sycl/source/detail/program_impl.cpp +++ b/sycl/source/detail/program_impl.cpp @@ -129,7 +129,7 @@ program_impl::program_impl(ContextImplPtr Context, "No InteropProgram/PiProgram defined with piextProgramFromNative"); // Translate the raw program handle into PI program. Plugin.call( - InteropProgram, MContext->getHandleRef(), &MProgram); + InteropProgram, MContext->getHandleRef(), &MProgram, false); } else Plugin.call(Program); diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index d42ded40e098d..50197bcb7dc39 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3834,6 +3834,7 @@ _ZN2cl4sycl6detail18convertChannelTypeE22_pi_image_channel_type _ZN2cl4sycl6detail18convertChannelTypeENS0_18image_channel_typeE _ZN2cl4sycl6detail18get_kernel_id_implENSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE _ZN2cl4sycl6detail18make_kernel_bundleEmRKNS0_7contextENS0_12bundle_stateENS0_7backendE +_ZN2cl4sycl6detail18make_kernel_bundleEmRKNS0_7contextEbNS0_12bundle_stateENS0_7backendE _ZN2cl4sycl6detail18stringifyErrorCodeEi _ZN2cl4sycl6detail19convertChannelOrderE23_pi_image_channel_order _ZN2cl4sycl6detail19convertChannelOrderENS0_19image_channel_orderE @@ -3914,8 +3915,8 @@ _ZN2cl4sycl7handler10mem_adviseEPKvmi _ZN2cl4sycl7handler10processArgEPvRKNS0_6detail19kernel_param_kind_tEimRmb _ZN2cl4sycl7handler10processArgEPvRKNS0_6detail19kernel_param_kind_tEimRmbb _ZN2cl4sycl7handler13getKernelNameB5cxx11Ev -_ZN2cl4sycl7handler18ext_oneapi_barrierERKSt6vectorINS0_5eventESaIS3_EE _ZN2cl4sycl7handler18RangeRoundingTraceEv +_ZN2cl4sycl7handler18ext_oneapi_barrierERKSt6vectorINS0_5eventESaIS3_EE _ZN2cl4sycl7handler18extractArgsAndReqsEv _ZN2cl4sycl7handler20DisableRangeRoundingEv _ZN2cl4sycl7handler20associateWithHandlerEPNS0_6detail16AccessorBaseHostENS0_6access6targetE