[SYCL] Fix kernel bundles don't really carry kernel IDs (#5121)

denis-kabanov · web-flow · commit 13443a9628df · 2021-12-22T11:28:05.000+03:00
Fix that interop kernel bundles don't really carry kernel IDs for contained kernels. This change adds In enqueueImpKernel another condition to taking the path using the associated kernel_bundle, namely that interop bundles can't be used in the first branch. (This is 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.)
diff --git 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<device> &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<device> &Devs);
+
+/// make_kernel may need an empty interop kernel bundle. This function supplies
+/// this.
+template <bundle_state State>
+kernel_bundle<State> 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<sycl::kernel_bundle<State>>(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
@@ -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<bundle_state::executable>(
-                         TargetContext, std::vector<kernel_id>{}),
-                     NativeHandle, false, Backend);
+  return make_kernel(
+      TargetContext,
+      get_empty_interop_kernel_bundle<bundle_state::executable>(TargetContext),
+      NativeHandle, false, Backend);
 }
 
 } // namespace detail
diff --git 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<device> Devs,
-                     device_image_plain &DevImage)
+  // Interop constructor used by make_kernel
+  kernel_bundle_impl(context Ctx, std::vector<device> 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<device> 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
@@ -1922,15 +1922,16 @@ cl_int enqueueImpKernel(
   std::shared_ptr<kernel_impl> SyclKernelImpl;
   std::shared_ptr<device_image_impl> DeviceImageImpl;
 
-  // Use kernel_bundle is available
-  if (KernelBundleImplPtr) {
-
-    std::shared_ptr<kernel_id_impl> KernelIDImpl =
-        std::make_shared<kernel_id_impl>(KernelName);
-
-    kernel SyclKernel = KernelBundleImplPtr->get_kernel(
-        detail::createSyclObjFromImpl<kernel_id>(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
@@ -138,6 +138,12 @@ get_kernel_bundle_impl(const context &Ctx, const std::vector<device> &Devs,
                                                       State);
 }
 
+detail::KernelBundleImplPtr
+get_empty_interop_kernel_bundle_impl(const context &Ctx,
+                                     const std::vector<device> &Devs) {
+  return std::make_shared<detail::kernel_bundle_impl>(Ctx, Devs);
+}
+
 std::shared_ptr<detail::kernel_bundle_impl>
 join_impl(const std::vector<detail::KernelBundleImplPtr> &Bundles) {
   return std::make_shared<detail::kernel_bundle_impl>(Bundles);
diff --git 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
@@ -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
