diff --git a/sycl/include/CL/sycl/kernel_bundle.hpp b/sycl/include/CL/sycl/kernel_bundle.hpp index 854b76aef2b5e..81226a7af7dad 100644 --- a/sycl/include/CL/sycl/kernel_bundle.hpp +++ b/sycl/include/CL/sycl/kernel_bundle.hpp @@ -471,6 +471,21 @@ using DevImgSelectorImpl = __SYCL_EXPORT detail::KernelBundleImplPtr get_kernel_bundle_impl(const context &Ctx, const std::vector &Devs, bundle_state State, const DevImgSelectorImpl &Selector); + +// Internal non-template versions of get_empty_interop_kernel_bundle API which +// is used by public onces +__SYCL_EXPORT detail::KernelBundleImplPtr +get_empty_interop_kernel_bundle_impl(const context &Ctx, + const std::vector &Devs); + +/// make_kernel may need an empty interop kernel bundle. This function supplies +/// this. +template +kernel_bundle get_empty_interop_kernel_bundle(const context &Ctx) { + detail::KernelBundleImplPtr Impl = + detail::get_empty_interop_kernel_bundle_impl(Ctx, Ctx.get_devices()); + return detail::createSyclObjFromImpl>(Impl); +} } // namespace detail /// A kernel bundle in state State which contains all of the device images for diff --git a/sycl/source/backend.cpp b/sycl/source/backend.cpp index a2a6a7a1561a4..691bdbf0ab91c 100644 --- a/sycl/source/backend.cpp +++ b/sycl/source/backend.cpp @@ -246,10 +246,10 @@ kernel make_kernel(const context &TargetContext, kernel make_kernel(pi_native_handle NativeHandle, const context &TargetContext, backend Backend) { - return make_kernel(TargetContext, - get_kernel_bundle( - TargetContext, std::vector{}), - NativeHandle, false, Backend); + return make_kernel( + TargetContext, + get_empty_interop_kernel_bundle(TargetContext), + NativeHandle, false, Backend); } } // namespace detail diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index b70910c88ca5b..fa4a336aa8a3f 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -87,19 +87,24 @@ class kernel_bundle_impl { MContext, MDevices, State); } - // Interop constructor - kernel_bundle_impl(context Ctx, std::vector Devs, - device_image_plain &DevImage) + // Interop constructor used by make_kernel + kernel_bundle_impl(context Ctx, std::vector Devs) : MContext(Ctx), MDevices(Devs) { if (!checkAllDevicesAreInContext(Devs, Ctx)) throw sycl::exception( make_error_code(errc::invalid), "Not all devices are associated with the context or " "vector of devices is empty"); - MDeviceImages.push_back(DevImage); MIsInterop = true; } + // Interop constructor + kernel_bundle_impl(context Ctx, std::vector Devs, + device_image_plain &DevImage) + : kernel_bundle_impl(Ctx, Devs) { + MDeviceImages.push_back(DevImage); + } + // Matches sycl::build and sycl::compile // Have one constructor because sycl::build and sycl::compile have the same // signature @@ -476,6 +481,9 @@ class kernel_bundle_impl { size_t size() const noexcept { return MDeviceImages.size(); } bundle_state get_bundle_state() const { + // Interop kernel-bundles are always in executable state + if (MIsInterop) + return bundle_state::executable; // All device images are expected to have the same state return MDeviceImages.empty() ? bundle_state::input diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp index 7bc32896e0280..c5ff9eb30dfaa 100644 --- a/sycl/source/detail/scheduler/commands.cpp +++ b/sycl/source/detail/scheduler/commands.cpp @@ -1922,15 +1922,16 @@ cl_int enqueueImpKernel( std::shared_ptr SyclKernelImpl; std::shared_ptr DeviceImageImpl; - // Use kernel_bundle is available - if (KernelBundleImplPtr) { - - std::shared_ptr KernelIDImpl = - std::make_shared(KernelName); - - kernel SyclKernel = KernelBundleImplPtr->get_kernel( - detail::createSyclObjFromImpl(KernelIDImpl), - KernelBundleImplPtr); + // Use kernel_bundle if available unless it is interop. + // Interop bundles can't be used in the first branch, because the kernels + // in interop kernel bundles (if any) do not have kernel_id + // and can therefore not be looked up, but since they are self-contained + // they can simply be launched directly. + if (KernelBundleImplPtr && !KernelBundleImplPtr->isInterop()) { + kernel_id KernelID = + detail::ProgramManager::getInstance().getSYCLKernelID(KernelName); + kernel SyclKernel = + KernelBundleImplPtr->get_kernel(KernelID, KernelBundleImplPtr); SyclKernelImpl = detail::getSyclObjImpl(SyclKernel); diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index 1e52a424e5088..a0aeb84ddb3ff 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -138,6 +138,12 @@ get_kernel_bundle_impl(const context &Ctx, const std::vector &Devs, State); } +detail::KernelBundleImplPtr +get_empty_interop_kernel_bundle_impl(const context &Ctx, + const std::vector &Devs) { + return std::make_shared(Ctx, Devs); +} + std::shared_ptr join_impl(const std::vector &Bundles) { return std::make_shared(Bundles); diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index b1ebca155551f..faa69d0680b37 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3891,6 +3891,7 @@ _ZN2cl4sycl6detail2pi9assertionEbPKc _ZN2cl4sycl6detail2pi9getPluginILNS0_7backendE1EEERKNS1_6pluginEv _ZN2cl4sycl6detail2pi9getPluginILNS0_7backendE2EEERKNS1_6pluginEv _ZN2cl4sycl6detail2pi9getPluginILNS0_7backendE5EEERKNS1_6pluginEv +_ZN2cl4sycl6detail36get_empty_interop_kernel_bundle_implERKNS0_7contextERKSt6vectorINS0_6deviceESaIS6_EE _ZN2cl4sycl6detail6OSUtil10getDirNameB5cxx11EPKc _ZN2cl4sycl6detail6OSUtil11alignedFreeEPv _ZN2cl4sycl6detail6OSUtil12alignedAllocEmm @@ -4243,7 +4244,6 @@ _ZNK2cl4sycl6kernel11get_backendEv _ZNK2cl4sycl6kernel11get_contextEv _ZNK2cl4sycl6kernel11get_programEv _ZNK2cl4sycl6kernel13getNativeImplEv -_ZNK2cl4sycl6kernel9getNativeEv _ZNK2cl4sycl6kernel17get_kernel_bundleEv _ZNK2cl4sycl6kernel18get_sub_group_infoILNS0_4info16kernel_sub_groupE16650EEENS3_12param_traitsIS4_XT_EE11return_typeERKNS0_6deviceE _ZNK2cl4sycl6kernel18get_sub_group_infoILNS0_4info16kernel_sub_groupE4537EEENS3_12param_traitsIS4_XT_EE11return_typeERKNS0_6deviceE @@ -4272,6 +4272,7 @@ _ZNK2cl4sycl6kernel8get_infoILNS0_4info6kernelE4498EEENS3_12param_traitsIS4_XT_E _ZNK2cl4sycl6kernel8get_infoILNS0_4info6kernelE4499EEENS3_12param_traitsIS4_XT_EE11return_typeEv _ZNK2cl4sycl6kernel8get_infoILNS0_4info6kernelE4500EEENS3_12param_traitsIS4_XT_EE11return_typeEv _ZNK2cl4sycl6kernel8get_infoILNS0_4info6kernelE4501EEENS3_12param_traitsIS4_XT_EE11return_typeEv +_ZNK2cl4sycl6kernel9getNativeEv _ZNK2cl4sycl6stream22get_max_statement_sizeEv _ZNK2cl4sycl6stream8get_sizeEv _ZNK2cl4sycl6streameqERKS1_ diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index 1d2507ee78dd5..c90dc3fa058f3 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -1873,7 +1873,6 @@ ?fill_usm@MemoryManager@detail@sycl@cl@@SAXPEAXV?$shared_ptr@Vqueue_impl@detail@sycl@cl@@@std@@_KHV?$vector@PEAU_pi_event@@V?$allocator@PEAU_pi_event@@@std@@@6@AEAPEAU_pi_event@@@Z ?fill_usm@MemoryManager@detail@sycl@cl@@SAXPEAXV?$shared_ptr@Vqueue_impl@detail@sycl@cl@@@std@@_KHV?$vector@PEAU_pi_event@@V?$allocator@PEAU_pi_event@@@std@@@6@PEAPEAU_pi_event@@@Z ?finalize@handler@sycl@cl@@AEAA?AVevent@23@XZ -?getESIMDDeviceInterface@detail@sycl@cl@@YAPEAUESIMDDeviceInterface@123@XZ ?find_device_intersection@detail@sycl@cl@@YA?AV?$vector@Vdevice@sycl@cl@@V?$allocator@Vdevice@sycl@cl@@@std@@@std@@AEBV?$vector@V?$kernel_bundle@$00@sycl@cl@@V?$allocator@V?$kernel_bundle@$00@sycl@cl@@@std@@@5@@Z ?floor@__host_std@cl@@YA?AV?$vec@M$00@sycl@2@V342@@Z ?floor@__host_std@cl@@YA?AV?$vec@M$01@sycl@2@V342@@Z @@ -2094,6 +2093,7 @@ ?getDevices@?$image_impl@$01@detail@sycl@cl@@AEAA?AV?$vector@Vdevice@sycl@cl@@V?$allocator@Vdevice@sycl@cl@@@std@@@std@@V?$shared_ptr@Vcontext_impl@detail@sycl@cl@@@6@@Z ?getDevices@?$image_impl@$02@detail@sycl@cl@@AEAA?AV?$vector@Vdevice@sycl@cl@@V?$allocator@Vdevice@sycl@cl@@@std@@@std@@V?$shared_ptr@Vcontext_impl@detail@sycl@cl@@@6@@Z ?getDirName@OSUtil@detail@sycl@cl@@SA?AV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@PEBD@Z +?getESIMDDeviceInterface@detail@sycl@cl@@YAPEAUESIMDDeviceInterface@123@XZ ?getElementSize@?$image_impl@$00@detail@sycl@cl@@QEBA_KXZ ?getElementSize@?$image_impl@$01@detail@sycl@cl@@QEBA_KXZ ?getElementSize@?$image_impl@$02@detail@sycl@cl@@QEBA_KXZ @@ -2116,13 +2116,13 @@ ?getNative@device@sycl@cl@@AEBA_KXZ ?getNative@device_image_plain@detail@sycl@cl@@QEBA_KXZ ?getNative@event@sycl@cl@@AEBA_KXZ +?getNative@kernel@sycl@cl@@AEBA_KXZ ?getNative@platform@sycl@cl@@AEBA_KXZ ?getNative@program@sycl@cl@@AEBA_KXZ ?getNative@queue@sycl@cl@@AEBA_KXZ ?getNativeContext@interop_handle@sycl@cl@@AEBA_KXZ ?getNativeDevice@interop_handle@sycl@cl@@AEBA_KXZ ?getNativeImpl@kernel@sycl@cl@@AEBA_KXZ -?getNative@kernel@sycl@cl@@AEBA_KXZ ?getNativeMem@interop_handle@sycl@cl@@AEBA_KPEAVAccessorImplHost@detail@23@@Z ?getNativeQueue@interop_handle@sycl@cl@@AEBA_KXZ ?getOSMemSize@OSUtil@detail@sycl@cl@@SA_KXZ @@ -2180,6 +2180,7 @@ ?get_devices@kernel_bundle_plain@detail@sycl@cl@@QEBA?AV?$vector@Vdevice@sycl@cl@@V?$allocator@Vdevice@sycl@cl@@@std@@@std@@XZ ?get_devices@platform@sycl@cl@@QEBA?AV?$vector@Vdevice@sycl@cl@@V?$allocator@Vdevice@sycl@cl@@@std@@@std@@W4device_type@info@23@@Z ?get_devices@program@sycl@cl@@QEBA?AV?$vector@Vdevice@sycl@cl@@V?$allocator@Vdevice@sycl@cl@@@std@@@std@@XZ +?get_empty_interop_kernel_bundle_impl@detail@sycl@cl@@YA?AV?$shared_ptr@Vkernel_bundle_impl@detail@sycl@cl@@@std@@AEBVcontext@23@AEBV?$vector@Vdevice@sycl@cl@@V?$allocator@Vdevice@sycl@cl@@@std@@@5@@Z ?get_filtering_mode@sampler@sycl@cl@@QEBA?AW4filtering_mode@23@XZ ?get_filtering_mode@sampler_impl@detail@sycl@cl@@QEBA?AW4filtering_mode@34@XZ ?get_flags@stream@sycl@cl@@AEBAIXZ