[SYCL] Remove duplicate devices on submission to kernel_bundle API functions (#4790)

When submitting devices to compile, section 4.11.7 of the SYCL 2020 spec states that the compilation is for each unique device, duplicate devices removed. Here we ensure that the devices sent to the backend are each unique.
Also, Level Zero doesn't support compilation for more than one device but it calls die() presently. Changing that as well.

Signed-off-by: Chris Perkins <[email protected]>

Author: cperkinsintel

sycl/include/CL/sycl/device.hpp

@@ -25,6 +25,7 @@ namespace sycl {
 class device_selector;
 namespace detail {
 class device_impl;
+auto getDeviceComparisonLambda();
 }
 
 /// The SYCL device class encapsulates a single SYCL device on which kernels
@@ -215,6 +216,8 @@ class __SYCL_EXPORT device {
 
   template <class T>
   friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
+
+  friend auto detail::getDeviceComparisonLambda();
 };
 
 } // namespace sycl

sycl/include/CL/sycl/kernel_bundle.hpp

@@ -19,6 +19,7 @@
 
 #include <cassert>
 #include <memory>
+#include <set>
 #include <vector>
 
 __SYCL_INLINE_NAMESPACE(cl) {
@@ -375,6 +376,21 @@ namespace detail {
 __SYCL_EXPORT detail::KernelBundleImplPtr
 get_kernel_bundle_impl(const context &Ctx, const std::vector<device> &Devs,
                        bundle_state State);
+
+inline auto getDeviceComparisonLambda() {
+  return [](device a, device b) { return a.getNative() != b.getNative(); };
+}
+
+inline const std::vector<device>
+removeDuplicateDevices(const std::vector<device> &Devs) {
+  auto compareDevices = getDeviceComparisonLambda();
+  std::set<device, decltype(compareDevices)> UniqueDeviceSet(
+      Devs.begin(), Devs.end(), compareDevices);
+  std::vector<device> UniqueDevices(UniqueDeviceSet.begin(),
+                                    UniqueDeviceSet.end());
+  return UniqueDevices;
+}
+
 } // namespace detail
 
 /// A kernel bundle in state State which contains all of the kernels in the
@@ -384,8 +400,10 @@ get_kernel_bundle_impl(const context &Ctx, const std::vector<device> &Devs,
 template <bundle_state State>
 kernel_bundle<State> get_kernel_bundle(const context &Ctx,
                                        const std::vector<device> &Devs) {
+  std::vector<device> UniqueDevices = detail::removeDuplicateDevices(Devs);
+
   detail::KernelBundleImplPtr Impl =
-      detail::get_kernel_bundle_impl(Ctx, Devs, State);
+      detail::get_kernel_bundle_impl(Ctx, UniqueDevices, State);
 
   return detail::createSyclObjFromImpl<kernel_bundle<State>>(Impl);
 }
@@ -417,8 +435,10 @@ template <bundle_state State>
 kernel_bundle<State>
 get_kernel_bundle(const context &Ctx, const std::vector<device> &Devs,
                   const std::vector<kernel_id> &KernelIDs) {
+  std::vector<device> UniqueDevices = detail::removeDuplicateDevices(Devs);
+
   detail::KernelBundleImplPtr Impl =
-      detail::get_kernel_bundle_impl(Ctx, Devs, KernelIDs, State);
+      detail::get_kernel_bundle_impl(Ctx, UniqueDevices, KernelIDs, State);
   return detail::createSyclObjFromImpl<kernel_bundle<State>>(Impl);
 }
 
@@ -459,14 +479,16 @@ template <bundle_state State, typename SelectorT>
 kernel_bundle<State> get_kernel_bundle(const context &Ctx,
                                        const std::vector<device> &Devs,
                                        SelectorT Selector) {
+  std::vector<device> UniqueDevices = detail::removeDuplicateDevices(Devs);
+
   detail::DevImgSelectorImpl SelectorWrapper =
       [Selector](const detail::DeviceImageImplPtr &DevImg) {
         return Selector(
             detail::createSyclObjFromImpl<sycl::device_image<State>>(DevImg));
       };
 
-  detail::KernelBundleImplPtr Impl =
-      detail::get_kernel_bundle_impl(Ctx, Devs, State, SelectorWrapper);
+  detail::KernelBundleImplPtr Impl = detail::get_kernel_bundle_impl(
+      Ctx, UniqueDevices, State, SelectorWrapper);
 
   return detail::createSyclObjFromImpl<sycl::kernel_bundle<State>>(Impl);
 }
@@ -589,8 +611,10 @@ compile_impl(const kernel_bundle<bundle_state::input> &InputBundle,
 inline kernel_bundle<bundle_state::object>
 compile(const kernel_bundle<bundle_state::input> &InputBundle,
         const std::vector<device> &Devs, const property_list &PropList = {}) {
+  std::vector<device> UniqueDevices = detail::removeDuplicateDevices(Devs);
+
   detail::KernelBundleImplPtr Impl =
-      detail::compile_impl(InputBundle, Devs, PropList);
+      detail::compile_impl(InputBundle, UniqueDevices, PropList);
   return detail::createSyclObjFromImpl<
       kernel_bundle<sycl::bundle_state::object>>(Impl);
 }
@@ -622,8 +646,10 @@ link_impl(const std::vector<kernel_bundle<bundle_state::object>> &ObjectBundles,
 inline kernel_bundle<bundle_state::executable>
 link(const std::vector<kernel_bundle<bundle_state::object>> &ObjectBundles,
      const std::vector<device> &Devs, const property_list &PropList = {}) {
+  std::vector<device> UniqueDevices = detail::removeDuplicateDevices(Devs);
+
   detail::KernelBundleImplPtr Impl =
-      detail::link_impl(ObjectBundles, Devs, PropList);
+      detail::link_impl(ObjectBundles, UniqueDevices, PropList);
   return detail::createSyclObjFromImpl<
       kernel_bundle<sycl::bundle_state::executable>>(Impl);
 }
@@ -667,8 +693,10 @@ build_impl(const kernel_bundle<bundle_state::input> &InputBundle,
 inline kernel_bundle<bundle_state::executable>
 build(const kernel_bundle<bundle_state::input> &InputBundle,
       const std::vector<device> &Devs, const property_list &PropList = {}) {
+  std::vector<device> UniqueDevices = detail::removeDuplicateDevices(Devs);
+
   detail::KernelBundleImplPtr Impl =
-      detail::build_impl(InputBundle, Devs, PropList);
+      detail::build_impl(InputBundle, UniqueDevices, PropList);
   return detail::createSyclObjFromImpl<
       kernel_bundle<sycl::bundle_state::executable>>(Impl);
 }

sycl/plugins/level_zero/pi_level_zero.cpp

@@ -3407,8 +3407,10 @@ pi_result piProgramCreateWithBinary(
   PI_ASSERT(Program, PI_INVALID_PROGRAM);
 
   // For now we support only one device.
-  if (NumDevices != 1)
-    die("piProgramCreateWithBinary: level_zero supports only one device.");
+  if (NumDevices != 1) {
+    zePrint("piProgramCreateWithBinary: level_zero supports only one device.");
+    return PI_INVALID_VALUE;
+  }
   if (!Binaries[0] || !Lengths[0]) {
     if (BinaryStatus)
       *BinaryStatus = PI_INVALID_VALUE;
@@ -3605,8 +3607,10 @@ pi_result piProgramLink(pi_context Context, pi_uint32 NumDevices,
 
   // We only support one device with Level Zero currently.
   pi_device Device = Context->Devices[0];
-  if (NumDevices != 1)
-    die("piProgramLink: level_zero supports only one device.");
+  if (NumDevices != 1) {
+    zePrint("piProgramLink: level_zero supports only one device.");
+    return PI_INVALID_VALUE;
+  }
 
   PI_ASSERT(DeviceList && DeviceList[0] == Device, PI_INVALID_DEVICE);
   PI_ASSERT(!PFnNotify && !UserData, PI_INVALID_VALUE);
@@ -3783,8 +3787,10 @@ static pi_result compileOrBuild(pi_program Program, pi_uint32 NumDevices,
   // We only support build to one device with Level Zero now.
   // TODO: we should eventually build to the possibly multiple root
   // devices in the context.
-  if (NumDevices != 1)
-    die("compileOrBuild: level_zero supports only one device.");
+  if (NumDevices != 1) {
+    zePrint("compileOrBuild: level_zero supports only one device.");
+    return PI_INVALID_VALUE;
+  }
 
   PI_ASSERT(DeviceList, PI_INVALID_DEVICE);