Adjusting for comments

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

Author: steffenlarsen

sycl/include/CL/sycl/handler.hpp

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

sycl/source/detail/kernel_bundle_impl.hpp

@@ -79,8 +79,7 @@ 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)),
-        MInitialState(State) {
+      : MContext(std::move(Ctx)), MDevices(std::move(Devs)), MState(State) {
 
     common_ctor_checks(State);
 
@@ -90,7 +89,7 @@ class kernel_bundle_impl {
 
   // Interop constructor used by make_kernel
   kernel_bundle_impl(context Ctx, std::vector<device> Devs)
-      : MContext(Ctx), MDevices(Devs), MInitialState(bundle_state::executable) {
+      : MContext(Ctx), MDevices(Devs), MState(bundle_state::executable) {
     if (!checkAllDevicesAreInContext(Devs, Ctx))
       throw sycl::exception(
           make_error_code(errc::invalid),
@@ -113,7 +112,7 @@ class kernel_bundle_impl {
                      std::vector<device> Devs, const property_list &PropList,
                      bundle_state TargetState)
       : MContext(InputBundle.get_context()), MDevices(std::move(Devs)),
-        MInitialState(TargetState) {
+        MState(TargetState) {
 
     MSpecConstValues = getSyclObjImpl(InputBundle)->get_spec_const_map_ref();
 
@@ -163,7 +162,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)), MInitialState(bundle_state::executable) {
+      : MDevices(std::move(Devs)), MState(bundle_state::executable) {
 
     if (MDevices.empty())
       throw sycl::exception(make_error_code(errc::invalid),
@@ -243,8 +242,7 @@ 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)),
-        MInitialState(State) {
+      : MContext(std::move(Ctx)), MDevices(std::move(Devs)), MState(State) {
 
     // TODO: Add a check that all kernel ids are compatible with at least one
     // device in Devs
@@ -256,8 +254,7 @@ 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)),
-        MInitialState(State) {
+      : MContext(std::move(Ctx)), MDevices(std::move(Devs)), MState(State) {
 
     common_ctor_checks(State);
 
@@ -268,7 +265,7 @@ class kernel_bundle_impl {
   // C'tor matches sycl::join API
   kernel_bundle_impl(const std::vector<detail::KernelBundleImplPtr> &Bundles,
                      bundle_state State)
-      : MInitialState(State) {
+      : MState(State) {
     if (Bundles.empty())
       return;
 
@@ -492,7 +489,7 @@ class kernel_bundle_impl {
       return bundle_state::executable;
     // All device images are expected to have the same state
     return MDeviceImages.empty()
-               ? MInitialState
+               ? MState
                : detail::getSyclObjImpl(MDeviceImages[0])->get_state();
   }
 
@@ -502,38 +499,31 @@ class kernel_bundle_impl {
 
   bool isInterop() const { return MIsInterop; }
 
-  void add_kernel(const kernel_id &KernelID, const device &Dev) {
+  bool add_kernel(const kernel_id &KernelID, const device &Dev) {
     // Skip if kernel is already there
     if (has_kernel(KernelID, Dev))
-      return;
+      return true;
 
     // 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.");
+    // No images found so we report as not inserted
+    if (NewDevImgs.empty())
+      return false;
 
     // Propagate already set specialization constants to the new images
-    for (device_image_plain DevImg : NewDevImgs)
+    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());
+    return true;
   }
 
 private:
@@ -544,7 +534,7 @@ class kernel_bundle_impl {
   // from any device image.
   SpecConstMapT MSpecConstValues;
   bool MIsInterop = false;
-  bundle_state MInitialState;
+  bundle_state MState;
 };
 
 } // namespace detail

sycl/source/handler.cpp

@@ -126,66 +126,6 @@ handler::getOrInsertHandlerKernelBundle(bool Insert) const {
   return KernelBundleImpPtr;
 }
 
-// Returns a shared_ptr to kernel_bundle stored in the extended members vector.
-//
-// If there is no kernel_bundle created:
-// returns newly created kernel_bundle with KernelId if Insert is true
-// returns shared_ptr(nullptr) if Insert is false
-//
-// If there already existed an implicitly created kernel_bundle, the kernel is
-// inserted into that bundle.
-std::shared_ptr<detail::kernel_bundle_impl>
-handler::getOrInsertHandlerKernelBundle(bool Insert,
-                                        const kernel_id &KernelId) const {
-
-  std::lock_guard<std::mutex> Lock(
-      detail::GlobalHandler::instance().getHandlerExtendedMembersMutex());
-
-  assert(!MSharedPtrStorage.empty());
-
-  std::shared_ptr<std::vector<detail::ExtendedMemberT>> ExtendedMembersVec =
-      detail::convertToExtendedMembers(MSharedPtrStorage[0]);
-  // Look for the kernel bundle in extended members
-  std::shared_ptr<detail::kernel_bundle_impl> KernelBundleImpPtr;
-  for (const detail::ExtendedMemberT &EMember : *ExtendedMembersVec)
-    if (detail::ExtendedMembersType::HANDLER_KERNEL_BUNDLE == EMember.MType) {
-      KernelBundleImpPtr =
-          std::static_pointer_cast<detail::kernel_bundle_impl>(EMember.MData);
-      break;
-    }
-
-  // No kernel bundle yet, create one
-  if (!KernelBundleImpPtr) {
-    if (Insert) {
-      KernelBundleImpPtr =
-          detail::getSyclObjImpl(get_kernel_bundle<bundle_state::input>(
-              MQueue->get_context(), {MQueue->get_device()}, {KernelId}));
-      if (KernelBundleImpPtr->empty()) {
-        KernelBundleImpPtr =
-            detail::getSyclObjImpl(get_kernel_bundle<bundle_state::executable>(
-                MQueue->get_context(), {MQueue->get_device()}, {KernelId}));
-      }
-
-      detail::ExtendedMemberT EMember = {
-          detail::ExtendedMembersType::HANDLER_KERNEL_BUNDLE,
-          KernelBundleImpPtr};
-      ExtendedMembersVec->push_back(EMember);
-    }
-    return KernelBundleImpPtr;
-  }
-
-  auto HandlerImplMember = (*ExtendedMembersVec)[0];
-  assert(detail::ExtendedMembersType::HANDLER_IMPL == HandlerImplMember.MType);
-  auto HandlerImpl =
-      std::static_pointer_cast<detail::handler_impl>(HandlerImplMember.MData);
-
-  // Kernel bundles set explicitly by the user must not be filtered
-  if (!HandlerImpl->isStateExplicitKernelBundle())
-    KernelBundleImpPtr->add_kernel(KernelId, MQueue->get_device());
-
-  return KernelBundleImpPtr;
-}
-
 // Sets kernel bundle to the provided one. Either replaces existing one or
 // create a new entry in the extended members vector.
 void handler::setHandlerKernelBundle(
@@ -232,7 +172,26 @@ event handler::finalize() {
           !getHandlerImpl()->isStateExplicitKernelBundle()) {
         kernel_id KernelID =
             detail::ProgramManager::getInstance().getSYCLKernelID(MKernelName);
-        KernelBundleImpPtr->add_kernel(KernelID, MQueue->get_device());
+        bool KernelInserted =
+            KernelBundleImpPtr->add_kernel(KernelID, MQueue->get_device());
+        // If kernel was not inserted and the bundle is in input mode we try
+        // building it and trying to find the kernel in executable mode
+        if (!KernelInserted &&
+            KernelBundleImpPtr->get_bundle_state() == bundle_state::input) {
+          auto KernelBundle =
+              detail::createSyclObjFromImpl<kernel_bundle<bundle_state::input>>(
+                  KernelBundleImpPtr);
+          kernel_bundle<bundle_state::executable> ExecKernelBundle =
+              build(KernelBundle);
+          KernelBundleImpPtr = detail::getSyclObjImpl(ExecKernelBundle);
+          setHandlerKernelBundle(KernelBundleImpPtr);
+          KernelInserted =
+            KernelBundleImpPtr->add_kernel(KernelID, MQueue->get_device());
+        }
+        // If the kernel was not found in executable mode we throw an exception
+        if (!KernelInserted)
+          throw runtime_error("Failed to add kernel to kernel bundle.",
+                              PI_INVALID_KERNEL_NAME);
       }
 
       switch (KernelBundleImpPtr->get_bundle_state()) {

sycl/test/abi/sycl_symbols_linux.dump

@@ -4321,7 +4321,6 @@ _ZNK2cl4sycl7context9getNativeEv
 _ZNK2cl4sycl7handler14getHandlerImplEv
 _ZNK2cl4sycl7handler27isStateExplicitKernelBundleEv
 _ZNK2cl4sycl7handler30getOrInsertHandlerKernelBundleEb
-_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

@@ -2132,7 +2132,6 @@
 ?getOSModuleHandle@OSUtil@detail@sycl@cl@@SA_JPEBX@Z
 ?getOrCreateSampler@sampler_impl@detail@sycl@cl@@QEAAPEAU_pi_sampler@@AEBVcontext@34@@Z
 ?getOrInsertHandlerKernelBundle@handler@sycl@cl@@AEBA?AV?$shared_ptr@Vkernel_bundle_impl@detail@sycl@cl@@@std@@_N@Z
-?getOrInsertHandlerKernelBundle@handler@sycl@cl@@AEBA?AV?$shared_ptr@Vkernel_bundle_impl@detail@sycl@cl@@@std@@_NAEBVkernel_id@23@@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
 ?getPixelCoordNearestFiltMode@detail@sycl@cl@@YA?AV?$vec@H$03@23@V?$vec@M$03@23@W4addressing_mode@23@V?$range@$02@23@@Z