Change of strategy

This commit changes the strategy from having implicitly created kernel
bundles include all device images and then filter later, to have the
implicitly created kernel bundles have no device images initially and
then add the images once required kernels are known.

Signed-off-by: Steffen Larsen <[email protected]>

Author: steffenlarsen

sycl/include/CL/sycl/handler.hpp

@@ -730,8 +730,8 @@ class __SYCL_EXPORT handler {
     // If the kernel lambda is callable with a kernel_handler argument, manifest
     // the associated kernel handler.
     if (IsCallableWithKernelHandler) {
-      getOrInsertFilteredHandlerKernelBundle(/*Insert=*/true,
-                                             {get_kernel_id<KernelName>()});
+      getOrInsertHandlerKernelBundle(/*Insert=*/true,
+                                     get_kernel_id<KernelName>());
     }
   }
 
@@ -1274,7 +1274,7 @@ class __SYCL_EXPORT handler {
   getOrInsertHandlerKernelBundle(bool Insert) const;
 
   std::shared_ptr<detail::kernel_bundle_impl>
-  getOrInsertFilteredHandlerKernelBundle(bool Insert, kernel_id KernelId) const;
+  getOrInsertHandlerKernelBundle(bool Insert, const kernel_id &KernelId) const;
 
   void setHandlerKernelBundle(
       const std::shared_ptr<detail::kernel_bundle_impl> &NewKernelBundleImpPtr);

sycl/include/CL/sycl/kernel_bundle.hpp

@@ -588,8 +588,13 @@ template <typename KernelName> bool is_compatible(const device &Dev) {
 
 namespace detail {
 
+// TODO: This is no longer in use. Remove when ABI break is allowed.
 __SYCL_EXPORT std::shared_ptr<detail::kernel_bundle_impl>
 join_impl(const std::vector<detail::KernelBundleImplPtr> &Bundles);
+
+__SYCL_EXPORT std::shared_ptr<detail::kernel_bundle_impl>
+join_impl(const std::vector<detail::KernelBundleImplPtr> &Bundles,
+          bundle_state State);
 }
 
 /// \returns a new kernel bundle that represents the union of all the device
@@ -604,7 +609,7 @@ join(const std::vector<sycl::kernel_bundle<State>> &Bundles) {
     KernelBundleImpls.push_back(detail::getSyclObjImpl(Bundle));
 
   std::shared_ptr<detail::kernel_bundle_impl> Impl =
-      detail::join_impl(KernelBundleImpls);
+      detail::join_impl(KernelBundleImpls, State);
   return detail::createSyclObjFromImpl<kernel_bundle<State>>(Impl);
 }
 

sycl/source/detail/kernel_bundle_impl.hpp

@@ -79,7 +79,8 @@ class kernel_bundle_impl {
 
 public:
   kernel_bundle_impl(context Ctx, std::vector<device> Devs, bundle_state State)
-      : MContext(std::move(Ctx)), MDevices(std::move(Devs)) {
+      : MContext(std::move(Ctx)), MDevices(std::move(Devs)),
+        MInitialState(State) {
 
     common_ctor_checks(State);
 
@@ -89,7 +90,7 @@ class kernel_bundle_impl {
 
   // Interop constructor used by make_kernel
   kernel_bundle_impl(context Ctx, std::vector<device> Devs)
-      : MContext(Ctx), MDevices(Devs) {
+      : MContext(Ctx), MDevices(Devs), MInitialState(bundle_state::executable) {
     if (!checkAllDevicesAreInContext(Devs, Ctx))
       throw sycl::exception(
           make_error_code(errc::invalid),
@@ -111,7 +112,8 @@ class kernel_bundle_impl {
   kernel_bundle_impl(const kernel_bundle<bundle_state::input> &InputBundle,
                      std::vector<device> Devs, const property_list &PropList,
                      bundle_state TargetState)
-      : MContext(InputBundle.get_context()), MDevices(std::move(Devs)) {
+      : MContext(InputBundle.get_context()), MDevices(std::move(Devs)),
+        MInitialState(TargetState) {
 
     MSpecConstValues = getSyclObjImpl(InputBundle)->get_spec_const_map_ref();
 
@@ -161,7 +163,7 @@ class kernel_bundle_impl {
   kernel_bundle_impl(
       const std::vector<kernel_bundle<bundle_state::object>> &ObjectBundles,
       std::vector<device> Devs, const property_list &PropList)
-      : MDevices(std::move(Devs)) {
+      : MDevices(std::move(Devs)), MInitialState(bundle_state::executable) {
 
     if (MDevices.empty())
       throw sycl::exception(make_error_code(errc::invalid),
@@ -241,7 +243,8 @@ class kernel_bundle_impl {
   kernel_bundle_impl(context Ctx, std::vector<device> Devs,
                      const std::vector<kernel_id> &KernelIDs,
                      bundle_state State)
-      : MContext(std::move(Ctx)), MDevices(std::move(Devs)) {
+      : MContext(std::move(Ctx)), MDevices(std::move(Devs)),
+        MInitialState(State) {
 
     // TODO: Add a check that all kernel ids are compatible with at least one
     // device in Devs
@@ -253,7 +256,8 @@ class kernel_bundle_impl {
 
   kernel_bundle_impl(context Ctx, std::vector<device> Devs,
                      const DevImgSelectorImpl &Selector, bundle_state State)
-      : MContext(std::move(Ctx)), MDevices(std::move(Devs)) {
+      : MContext(std::move(Ctx)), MDevices(std::move(Devs)),
+        MInitialState(State) {
 
     common_ctor_checks(State);
 
@@ -262,7 +266,8 @@ class kernel_bundle_impl {
   }
 
   // C'tor matches sycl::join API
-  kernel_bundle_impl(const std::vector<detail::KernelBundleImplPtr> &Bundles) {
+  kernel_bundle_impl(const std::vector<detail::KernelBundleImplPtr> &Bundles,
+                     bundle_state State) : MInitialState(State) {
     if (Bundles.empty())
       return;
 
@@ -486,7 +491,7 @@ class kernel_bundle_impl {
       return bundle_state::executable;
     // All device images are expected to have the same state
     return MDeviceImages.empty()
-               ? bundle_state::input
+               ? MInitialState
                : detail::getSyclObjImpl(MDeviceImages[0])->get_state();
   }
 
@@ -496,17 +501,38 @@ class kernel_bundle_impl {
 
   bool isInterop() const { return MIsInterop; }
 
-  void filterImages(const kernel_id &KernelId) {
-    auto ImgHasKernelPred = [&KernelId](const device_image_plain &Img) {
-      return Img.has_kernel(KernelId);
-    };
-    const size_t NumRequiredDevImgs = std::count_if(
-        MDeviceImages.begin(), MDeviceImages.end(), ImgHasKernelPred);
-    std::vector<device_image_plain> FilteredDeviceImages;
-    FilteredDeviceImages.reserve(NumRequiredDevImgs);
-    std::copy_if(MDeviceImages.begin(), MDeviceImages.end(),
-                 std::back_inserter(FilteredDeviceImages), ImgHasKernelPred);
-    MDeviceImages = FilteredDeviceImages;
+  void add_kernel(const kernel_id &KernelID, const device &Dev) {
+    // Skip if kernel is already there
+    if (has_kernel(KernelID, Dev))
+      return;
+
+    // First try and get images in current bundle state
+    const bundle_state BundleState = get_bundle_state();
+    std::vector<device_image_plain> NewDevImgs =
+        detail::ProgramManager::getInstance().getSYCLDeviceImages(
+            MContext, {Dev}, {KernelID}, BundleState);
+
+    // If no images were found and the bundle is in input state we try and get
+    // the image in executable state and then bring the existing binaries into
+    // executable as well
+    if (NewDevImgs.empty() && BundleState == bundle_state::input) {
+      NewDevImgs = detail::ProgramManager::getInstance().getSYCLDeviceImages(
+          MContext, {Dev}, {KernelID}, bundle_state::executable);
+      detail::ProgramManager::getInstance().bringSYCLDeviceImagesToState(
+          MDeviceImages, bundle_state::executable);
+    }
+
+    assert(!NewDevImgs.empty() && "Device images for kernel was not found.");
+
+    // Propagate already set specialization constants to the new images
+    for (device_image_plain DevImg : NewDevImgs)
+      for (auto SpecConst : MSpecConstValues)
+        getSyclObjImpl(DevImg)->set_specialization_constant_raw_value(
+            SpecConst.first.c_str(), SpecConst.second.data());
+
+    // Add the images to the collection
+    MDeviceImages.insert(MDeviceImages.end(), NewDevImgs.begin(),
+                         NewDevImgs.end());
   }
 
 private:
@@ -517,6 +543,7 @@ class kernel_bundle_impl {
   // from any device image.
   SpecConstMapT MSpecConstValues;
   bool MIsInterop = false;
+  bundle_state MInitialState;
 };
 
 } // namespace detail

sycl/source/handler.cpp

@@ -113,12 +113,10 @@ handler::getOrInsertHandlerKernelBundle(bool Insert) const {
 
   // No kernel bundle yet, create one
   if (!KernelBundleImpPtr && Insert) {
-    KernelBundleImpPtr = detail::getSyclObjImpl(
-        get_kernel_bundle<bundle_state::input>(MQueue->get_context()));
-    if (KernelBundleImpPtr->empty()) {
-      KernelBundleImpPtr = detail::getSyclObjImpl(
-          get_kernel_bundle<bundle_state::executable>(MQueue->get_context()));
-    }
+    // Create an empty kernel bundle to add kernels to later
+    KernelBundleImpPtr =
+        detail::getSyclObjImpl(get_kernel_bundle<bundle_state::input>(
+            MQueue->get_context(), {MQueue->get_device()}, {}));
 
     detail::ExtendedMemberT EMember = {
         detail::ExtendedMembersType::HANDLER_KERNEL_BUNDLE, KernelBundleImpPtr};
@@ -134,12 +132,11 @@ handler::getOrInsertHandlerKernelBundle(bool Insert) const {
 // returns newly created kernel_bundle with KernelId if Insert is true
 // returns shared_ptr(nullptr) if Insert is false
 //
-// If there already existed a kernel_bundle, the underlying device images are
-// filtered such that only the ones containing KernelId remain in the
-// kernel_bundle.
+// If there already existed an implicitly created kernel_bundle, the kernel is
+// inserted into that bundle.
 std::shared_ptr<detail::kernel_bundle_impl>
-handler::getOrInsertFilteredHandlerKernelBundle(bool Insert,
-                                                kernel_id KernelId) const {
+handler::getOrInsertHandlerKernelBundle(bool Insert,
+                                        const kernel_id &KernelId) const {
 
   std::lock_guard<std::mutex> Lock(
       detail::GlobalHandler::instance().getHandlerExtendedMembersMutex());
@@ -162,11 +159,11 @@ handler::getOrInsertFilteredHandlerKernelBundle(bool Insert,
     if (Insert) {
       KernelBundleImpPtr =
           detail::getSyclObjImpl(get_kernel_bundle<bundle_state::input>(
-              MQueue->get_context(), {KernelId}));
+              MQueue->get_context(), {MQueue->get_device()}, {KernelId}));
       if (KernelBundleImpPtr->empty()) {
         KernelBundleImpPtr =
             detail::getSyclObjImpl(get_kernel_bundle<bundle_state::executable>(
-                MQueue->get_context(), {KernelId}));
+                MQueue->get_context(), {MQueue->get_device()}, {KernelId}));
       }
 
       detail::ExtendedMemberT EMember = {
@@ -184,7 +181,7 @@ handler::getOrInsertFilteredHandlerKernelBundle(bool Insert,
 
   // Kernel bundles set explicitly by the user must not be filtered
   if (!HandlerImpl->isStateExplicitKernelBundle())
-    KernelBundleImpPtr->filterImages(KernelId);
+    KernelBundleImpPtr->add_kernel(KernelId, MQueue->get_device());
 
   return KernelBundleImpPtr;
 }
@@ -230,6 +227,13 @@ event handler::finalize() {
     // If there were uses of set_specialization_constant build the kernel_bundle
     KernelBundleImpPtr = getOrInsertHandlerKernelBundle(/*Insert=*/false);
     if (KernelBundleImpPtr) {
+      // Make sure implicit kernel bundles has the kernel
+      if (!getHandlerImpl()->isStateExplicitKernelBundle()) {
+        kernel_id KernelID =
+            detail::ProgramManager::getInstance().getSYCLKernelID(MKernelName);
+        KernelBundleImpPtr->add_kernel(KernelID, MQueue->get_device());
+      }
+
       switch (KernelBundleImpPtr->get_bundle_state()) {
       case bundle_state::input: {
         // Underlying level expects kernel_bundle to be in executable state
@@ -679,6 +683,10 @@ void handler::verifyUsedKernelBundle(const std::string &KernelName) {
   if (!UsedKernelBundleImplPtr)
     return;
 
+  // Implicit kernel bundles are populated late so we ignore them
+  if (!getHandlerImpl()->isStateExplicitKernelBundle())
+    return;
+
   kernel_id KernelID = detail::get_kernel_id_impl(KernelName);
   device Dev = detail::getDeviceFromHandler(*this);
   if (!UsedKernelBundleImplPtr->has_kernel(KernelID, Dev))

sycl/source/kernel_bundle.cpp

@@ -146,7 +146,14 @@ get_empty_interop_kernel_bundle_impl(const context &Ctx,
 
 std::shared_ptr<detail::kernel_bundle_impl>
 join_impl(const std::vector<detail::KernelBundleImplPtr> &Bundles) {
-  return std::make_shared<detail::kernel_bundle_impl>(Bundles);
+  return std::make_shared<detail::kernel_bundle_impl>(Bundles,
+                                                      bundle_state::input);
+}
+
+std::shared_ptr<detail::kernel_bundle_impl>
+join_impl(const std::vector<detail::KernelBundleImplPtr> &Bundles,
+          bundle_state State) {
+  return std::make_shared<detail::kernel_bundle_impl>(Bundles, State);
 }
 
 bool has_kernel_bundle_impl(const context &Ctx, const std::vector<device> &Devs,

sycl/test/abi/sycl_symbols_linux.dump

@@ -3903,6 +3903,7 @@ _ZN2cl4sycl6detail6OSUtil16getCurrentDSODirB5cxx11Ev
 _ZN2cl4sycl6detail6OSUtil17getOSModuleHandleEPKv
 _ZN2cl4sycl6detail6OSUtil7makeDirEPKc
 _ZN2cl4sycl6detail9join_implERKSt6vectorISt10shared_ptrINS1_18kernel_bundle_implEESaIS5_EE
+_ZN2cl4sycl6detail9join_implERKSt6vectorISt10shared_ptrINS1_18kernel_bundle_implEESaIS5_EENS0_12bundle_stateE
 _ZN2cl4sycl6detail9link_implERKSt6vectorINS0_13kernel_bundleILNS0_12bundle_stateE1EEESaIS5_EERKS2_INS0_6deviceESaISA_EERKNS0_13property_listE
 _ZN2cl4sycl6device11get_devicesENS0_4info11device_typeE
 _ZN2cl4sycl6deviceC1EP13_cl_device_id
@@ -4320,7 +4321,7 @@ _ZNK2cl4sycl7context9getNativeEv
 _ZNK2cl4sycl7handler14getHandlerImplEv
 _ZNK2cl4sycl7handler27isStateExplicitKernelBundleEv
 _ZNK2cl4sycl7handler30getOrInsertHandlerKernelBundleEb
-_ZNK2cl4sycl7handler38getOrInsertFilteredHandlerKernelBundleEbNS0_9kernel_idE
+_ZNK2cl4sycl7handler30getOrInsertHandlerKernelBundleEbRKNS0_9kernel_idE
 _ZNK2cl4sycl7program10get_kernelENSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE
 _ZNK2cl4sycl7program10get_kernelENSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEb
 _ZNK2cl4sycl7program10has_kernelENSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE

sycl/test/abi/sycl_symbols_windows.dump

@@ -2131,7 +2131,6 @@
 ?getOSMemSize@OSUtil@detail@sycl@cl@@SA_KXZ
 ?getOSModuleHandle@OSUtil@detail@sycl@cl@@SA_JPEBX@Z
 ?getOrCreateSampler@sampler_impl@detail@sycl@cl@@QEAAPEAU_pi_sampler@@AEBVcontext@34@@Z
-?getOrInsertFilteredHandlerKernelBundle@handler@sycl@cl@@AEBA?AV?$shared_ptr@Vkernel_bundle_impl@detail@sycl@cl@@@std@@_NVkernel_id@23@@Z
 ?getOrInsertHandlerKernelBundle@handler@sycl@cl@@AEBA?AV?$shared_ptr@Vkernel_bundle_impl@detail@sycl@cl@@@std@@_N@Z
 ?getOrWaitEvents@detail@sycl@cl@@YA?AV?$vector@PEAU_pi_event@@V?$allocator@PEAU_pi_event@@@std@@@std@@V?$vector@Vevent@sycl@cl@@V?$allocator@Vevent@sycl@cl@@@std@@@5@V?$shared_ptr@Vcontext_impl@detail@sycl@cl@@@5@@Z
 ?getPixelCoordLinearFiltMode@detail@sycl@cl@@YA?AV?$vec@H$07@23@V?$vec@M$03@23@W4addressing_mode@23@V?$range@$02@23@AEAV523@@Z