Address comments

Author: romanovvlad

sycl/source/detail/kernel_bundle_impl.hpp

@@ -33,7 +33,7 @@ template <class T> struct LessByHash {
   }
 };
 
-static bool CheckAllDevicesAreInContext(const std::vector<device> &Devices,
+static bool checkAllDevicesAreInContext(const std::vector<device> &Devices,
                                         const context &Context) {
   const std::vector<device> &ContextDevices = Context.get_devices();
   return std::all_of(
@@ -43,7 +43,7 @@ static bool CheckAllDevicesAreInContext(const std::vector<device> &Devices,
       });
 }
 
-static bool CheckAllDevicesHaveAspect(const std::vector<device> &Devices,
+static bool checkAllDevicesHaveAspect(const std::vector<device> &Devices,
                                       aspect Aspect) {
   return std::all_of(Devices.begin(), Devices.end(),
                      [&Aspect](const device &Dev) { return Dev.has(Aspect); });
@@ -56,20 +56,20 @@ class kernel_bundle_impl {
 
   void common_ctor_checks(bundle_state State) {
     const bool AllDevicesInTheContext =
-        CheckAllDevicesAreInContext(MDevices, MContext);
+        checkAllDevicesAreInContext(MDevices, MContext);
     if (MDevices.empty() || !AllDevicesInTheContext)
       throw sycl::exception(
           make_error_code(errc::invalid),
           "Not all devices are associated with the context or "
           "vector of devices is empty");
 
     if (bundle_state::input == State &&
-        !CheckAllDevicesHaveAspect(MDevices, aspect::online_compiler))
+        !checkAllDevicesHaveAspect(MDevices, aspect::online_compiler))
       throw sycl::exception(make_error_code(errc::invalid),
                             "Not all devices have aspect::online_compiler");
 
     if (bundle_state::object == State &&
-        !CheckAllDevicesHaveAspect(MDevices, aspect::online_linker))
+        !checkAllDevicesHaveAspect(MDevices, aspect::online_linker))
       throw sycl::exception(make_error_code(errc::invalid),
                             "Not all devices have aspect::online_linker");
   }
@@ -155,6 +155,7 @@ class kernel_bundle_impl {
     // devices for any of the bundles in ObjectBundles
     const bool AllDevsAssociatedWithInputBundles = std::all_of(
         MDevices.begin(), MDevices.end(), [&ObjectBundles](const device &Dev) {
+          // Number of devices is expected to be small
           return std::all_of(
               ObjectBundles.begin(), ObjectBundles.end(),
               [&Dev](const kernel_bundle<bundle_state::object> &KernelBundle) {

sycl/source/kernel_bundle.cpp

@@ -117,17 +117,18 @@ join_impl(const std::vector<detail::KernelBundleImplPtr> &Bundles) {
 
 bool has_kernel_bundle_impl(const context &Ctx, const std::vector<device> &Devs,
                             bundle_state State) {
-  const bool AllDevicesInTheContext = CheckAllDevicesAreInContext(Devs, Ctx);
+  // Check that all requested devices are associated with the context
+  const bool AllDevicesInTheContext = checkAllDevicesAreInContext(Devs, Ctx);
   if (Devs.empty() || !AllDevicesInTheContext)
     throw sycl::exception(make_error_code(errc::invalid),
                           "Not all devices are associated with the context or "
                           "vector of devices is empty");
 
   if (bundle_state::input == State &&
-      !CheckAllDevicesHaveAspect(Devs, aspect::online_compiler))
+      !checkAllDevicesHaveAspect(Devs, aspect::online_compiler))
     return false;
   if (bundle_state::object == State &&
-      !CheckAllDevicesHaveAspect(Devs, aspect::online_linker))
+      !checkAllDevicesHaveAspect(Devs, aspect::online_linker))
     return false;
 
   const std::vector<device_image_plain> DeviceImages =
@@ -143,7 +144,8 @@ bool has_kernel_bundle_impl(const context &Ctx, const std::vector<device> &Devs,
 bool has_kernel_bundle_impl(const context &Ctx, const std::vector<device> &Devs,
                             const std::vector<kernel_id> &KernelIds,
                             bundle_state State) {
-  const bool AllDevicesInTheContext = CheckAllDevicesAreInContext(Devs, Ctx);
+  // Check that all requested devices are associated with the context
+  const bool AllDevicesInTheContext = checkAllDevicesAreInContext(Devs, Ctx);
 
   if (Devs.empty() || !AllDevicesInTheContext)
     throw sycl::exception(make_error_code(errc::invalid),
@@ -218,33 +220,18 @@ std::vector<sycl::device> find_device_intersection(
   // for all bundles
   std::vector<sycl::device> IntersectDevices;
   std::vector<unsigned int> DevsCounters;
+  std::map<device, unsigned int, LessByHash<device>> DevCounters;
   for (const sycl::kernel_bundle<bundle_state::object> &ObjectBundle :
        ObjectBundles)
-    // Increment counter in "DevsCounters" each time a device is seen
-    for (const sycl::device &Device : ObjectBundle.get_devices()) {
-      auto It =
-          std::find(IntersectDevices.begin(), IntersectDevices.end(), Device);
-      if (IntersectDevices.end() != It) {
-        assert((size_t)(std::distance(IntersectDevices.begin(), It) + 1) ==
-               DevsCounters.size());
-        ++DevsCounters[std::distance(IntersectDevices.begin(), It)];
-        continue;
-      }
-      IntersectDevices.push_back(Device);
-      DevsCounters.push_back(1);
-    }
-
-  // If for some device counter is less than ObjectBundles.size() it means some
-  // bundle doesn't have it - remove such a device from the final result
-  size_t NewSize = DevsCounters.size();
-  for (size_t Idx = 0; Idx < NewSize; ++Idx) {
-    if (ObjectBundles.size() == DevsCounters[Idx])
-      continue;
-
-    std::swap(IntersectDevices[Idx], IntersectDevices.back());
-    --NewSize;
-  }
-  IntersectDevices.resize(NewSize);
+    // Increment counter in "DevCounters" each time a device is seen
+    for (const sycl::device &Device : ObjectBundle.get_devices())
+      DevCounters[Device]++;
+
+  // If some device counter is less than ObjectBundles.size() then some bundle
+  // doesn't have it - do not add such a device to the final result
+  for (const std::pair<const device, unsigned int> &It : DevCounters)
+    if (ObjectBundles.size() == It.second)
+      IntersectDevices.push_back(It.first);
 
   return IntersectDevices;
 }