Merge from 'sycl' to 'sycl-web' (4 commits)

Author: iclsrc

sycl/source/detail/device_image_impl.hpp

@@ -393,9 +393,8 @@ class device_image_impl
                   const device &DeviceCand) const noexcept {
     // If the device is in the device list and the kernel ID is in the kernel
     // bundle, return true.
-    for (device_impl &Device : get_devices())
-      if (&Device == &*getSyclObjImpl(DeviceCand))
-        return has_kernel(KernelIDCand);
+    if (get_devices().contains(*getSyclObjImpl(DeviceCand)))
+      return has_kernel(KernelIDCand);
 
     // Otherwise, if the device candidate is a sub-device it is also valid if
     // its parent is valid.

sycl/source/detail/device_impl.hpp

@@ -2282,24 +2282,10 @@ class device_impl : public std::enable_shared_from_this<device_impl> {
 
 }; // class device_impl
 
-struct devices_deref_impl {
-  template <typename T> static device_impl &dereference(T &Elem) {
-    using Ty = std::decay_t<decltype(Elem)>;
-    if constexpr (std::is_same_v<Ty, device>) {
-      return *getSyclObjImpl(Elem);
-    } else if constexpr (std::is_same_v<Ty, device_impl>) {
-      return Elem;
-    } else {
-      return *Elem;
-    }
-  }
-};
-using devices_iterator =
-    variadic_iterator<devices_deref_impl, device,
-                      std::vector<std::shared_ptr<device_impl>>::const_iterator,
-                      std::vector<device>::const_iterator,
-                      std::vector<device_impl *>::const_iterator,
-                      device_impl *>;
+using devices_iterator = variadic_iterator<
+    device, std::vector<std::shared_ptr<device_impl>>::const_iterator,
+    std::vector<device>::const_iterator,
+    std::vector<device_impl *>::const_iterator, device_impl *>;
 
 class devices_range : public iterator_range<devices_iterator> {
 private:

sycl/source/detail/graph/graph_impl.cpp

@@ -376,19 +376,13 @@ std::set<node_impl *> graph_impl::getCGEdges(
   for (auto &Req : Requirements) {
     // Look through the graph for nodes which share this requirement
     for (node_impl &Node : nodes()) {
-      if (Node.hasRequirementDependency(Req)) {
-        bool ShouldAddDep = true;
-        // If any of this node's successors have this requirement then we skip
-        // adding the current node as a dependency.
-        for (node_impl &Succ : Node.successors()) {
-          if (Succ.hasRequirementDependency(Req)) {
-            ShouldAddDep = false;
-            break;
-          }
-        }
-        if (ShouldAddDep) {
-          UniqueDeps.insert(&Node);
-        }
+      if (Node.hasRequirementDependency(Req) &&
+          // If any of this node's successors have this requirement then we skip
+          // adding the current node as a dependency.
+          none_of(Node.successors(), [&](node_impl &Succ) {
+            return Succ.hasRequirementDependency(Req);
+          })) {
+        UniqueDeps.insert(&Node);
       }
     }
   }

sycl/source/detail/graph/node_impl.hpp

@@ -749,37 +749,13 @@ class node_impl : public std::enable_shared_from_this<node_impl> {
   }
 };
 
-struct nodes_deref_impl {
-  template <typename T> static node_impl &dereference(T &Elem) {
-    using Ty = std::decay_t<decltype(Elem)>;
-    if constexpr (std::is_same_v<Ty, std::weak_ptr<node_impl>>) {
-      // This assumes that weak_ptr doesn't actually manage lifetime and
-      // the object is guaranteed to be alive (which seems to be the
-      // assumption across all graph code).
-      return *Elem.lock();
-    } else if constexpr (std::is_same_v<Ty, node>) {
-      return *getSyclObjImpl(Elem);
-    } else {
-      return *Elem;
-    }
-  }
-};
-
 template <typename... ContainerTy>
 using nodes_iterator_impl =
-    variadic_iterator<nodes_deref_impl, node,
-                      typename ContainerTy::const_iterator...>;
+    variadic_iterator<node, typename ContainerTy::const_iterator...>;
 
 using nodes_iterator = nodes_iterator_impl<
     std::vector<std::shared_ptr<node_impl>>, std::vector<node_impl *>,
-    // Next one is temporary. It looks like `weak_ptr`s aren't
-    // used for the actual lifetime management and the objects are
-    // always guaranteed to be alive. Once the code is cleaned
-    // from `weak_ptr`s this alternative should be removed too.
-    std::vector<std::weak_ptr<node_impl>>,
-    //
     std::set<std::shared_ptr<node_impl>>, std::set<node_impl *>,
-    //
     std::list<node_impl *>, std::vector<node>>;
 
 class nodes_range : public iterator_range<nodes_iterator> {

sycl/source/detail/helpers.hpp

@@ -35,21 +35,19 @@ const RTDeviceBinaryImage *
 retrieveKernelBinary(queue_impl &Queue, KernelNameStrRefT KernelName,
                      CGExecKernel *CGKernel = nullptr);
 
-template <typename DereferenceImpl, typename SyclTy, typename... Iterators>
-class variadic_iterator {
+template <typename SyclTy, typename... Iterators> class variadic_iterator {
   using storage_iter = std::variant<Iterators...>;
 
   storage_iter It;
 
 public:
   using iterator_category = std::forward_iterator_tag;
   using difference_type = std::ptrdiff_t;
-  using reference = decltype(DereferenceImpl::dereference(
-      *std::declval<nth_type_t<0, Iterators...>>()));
-  using value_type = std::remove_reference_t<reference>;
+  using value_type = std::remove_reference_t<decltype(*getSyclObjImpl(
+      std::declval<SyclTy>()))>;
   using sycl_type = SyclTy;
   using pointer = value_type *;
-  static_assert(std::is_same_v<reference, value_type &>);
+  using reference = value_type &;
 
   variadic_iterator() = default;
   variadic_iterator(const variadic_iterator &) = default;
@@ -80,7 +78,16 @@ class variadic_iterator {
   decltype(auto) operator*() {
     return std::visit(
         [](auto &&It) -> decltype(auto) {
-          return DereferenceImpl::dereference(*It);
+          decltype(auto) Elem = *It;
+          using Ty = std::decay_t<decltype(Elem)>;
+          static_assert(!std::is_same_v<Ty, decltype(Elem)>);
+          if constexpr (std::is_same_v<Ty, sycl_type>) {
+            return *getSyclObjImpl(Elem);
+          } else if constexpr (std::is_same_v<Ty, value_type>) {
+            return Elem;
+          } else {
+            return *Elem;
+          }
         },
         It);
   }
@@ -172,6 +179,18 @@ template <typename iterator> class iterator_range {
   iterator Begin;
   iterator End;
   const size_t Size;
+
+  template <class Pred> friend bool all_of(iterator_range R, Pred &&P) {
+    return std::all_of(R.begin(), R.end(), std::forward<Pred>(P));
+  }
+
+  template <class Pred> friend bool any_of(iterator_range R, Pred &&P) {
+    return std::any_of(R.begin(), R.end(), std::forward<Pred>(P));
+  }
+
+  template <class Pred> friend bool none_of(iterator_range R, Pred &&P) {
+    return std::none_of(R.begin(), R.end(), std::forward<Pred>(P));
+  }
 };
 } // namespace detail
 } // namespace _V1

sycl/source/detail/program_manager/program_manager.cpp

@@ -118,33 +118,27 @@ static bool isDeviceBinaryTypeSupported(context_impl &ContextImpl,
   devices_range Devices = ContextImpl.getDevices();
 
   // Program type is SPIR-V, so we need a device compiler to do JIT.
-  for (device_impl &D : Devices) {
-    if (!D.get_info<info::device::is_compiler_available>())
-      return false;
-  }
+  if (!all_of(Devices, [](device_impl &D) {
+        return D.get_info<info::device::is_compiler_available>();
+      }))
+    return false;
 
   // OpenCL 2.1 and greater require clCreateProgramWithIL
   if (ContextBackend == backend::opencl) {
-    std::string ver = ContextImpl.get_info<info::context::platform>()
-                          .get_info<info::platform::version>();
+    std::string ver =
+        ContextImpl.getPlatformImpl().get_info<info::platform::version>();
     if (ver.find("OpenCL 1.0") == std::string::npos &&
         ver.find("OpenCL 1.1") == std::string::npos &&
         ver.find("OpenCL 1.2") == std::string::npos &&
         ver.find("OpenCL 2.0") == std::string::npos)
       return true;
   }
 
-  for (device_impl &D : Devices) {
-    // We need cl_khr_il_program extension to be present
-    // and we can call clCreateProgramWithILKHR using the extension
-    std::vector<std::string> Extensions =
-        D.get_info<info::device::extensions>();
-    if (Extensions.end() ==
-        std::find(Extensions.begin(), Extensions.end(), "cl_khr_il_program"))
-      return false;
-  }
-
-  return true;
+  // We need cl_khr_il_program extension to be present
+  // and we can call clCreateProgramWithILKHR using the extension
+  return all_of(Devices, [](device_impl &D) {
+    return D.has_extension("cl_khr_il_program");
+  });
 }
 
 // getFormatStr is used for debug-printing, so it may be unused.

sycl/source/detail/scheduler/graph_builder.cpp

@@ -674,11 +674,9 @@ static bool checkHostUnifiedMemory(context_impl *Ctx) {
   if (Ctx == nullptr)
     return true;
 
-  for (device_impl &Device : Ctx->getDevices()) {
-    if (!Device.get_info<info::device::host_unified_memory>())
-      return false;
-  }
-  return true;
+  return all_of(Ctx->getDevices(), [](device_impl &Device) {
+    return Device.get_info<info::device::host_unified_memory>();
+  });
 }
 
 // The function searches for the alloca command matching context and

sycl/test-e2e/FreeFunctionKernels/kernel_bundle.cpp

@@ -0,0 +1,85 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+#include <iostream>
+#include <sycl/detail/core.hpp>
+#include <sycl/ext/oneapi/get_kernel_info.hpp>
+#include <sycl/kernel_bundle.hpp>
+
+namespace syclext = sycl::ext::oneapi;
+namespace syclexp = sycl::ext::oneapi::experimental;
+
+static constexpr size_t NUM = 1024;
+static constexpr size_t WGSIZE = 16;
+static constexpr auto FFTestMark = "Free function Kernel Test:";
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<2>))
+void func_range(float start, float *ptr) {}
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::single_task_kernel))
+void func_single(float start, float *ptr) {}
+
+template <auto *Func> int test_free_function_kernel_id(sycl::context &ctxt) {
+  sycl::kernel_id id = syclexp::get_kernel_id<Func>();
+  auto exe_bndl =
+      syclexp::get_kernel_bundle<Func, sycl::bundle_state::executable>(ctxt);
+  const auto kernel_ids = exe_bndl.get_kernel_ids();
+  const bool res =
+      std::find(kernel_ids.begin(), kernel_ids.end(), id) == kernel_ids.end();
+  if (res)
+    std::cout << FFTestMark << "test_kernel_id failed: kernel id is not found"
+              << std::endl;
+  return res;
+}
+
+template <auto *Func>
+int test_kernel_bundle_ctxt(sycl::context &ctxt, std::string_view fname) {
+  sycl::kernel_id id = syclexp::get_kernel_id<Func>();
+  auto exe_bndl =
+      syclexp::get_kernel_bundle<Func, sycl::bundle_state::executable>(ctxt);
+  const bool res =
+      exe_bndl.has_kernel(id) &&
+      exe_bndl.get_kernel(id)
+              .template get_info<sycl::info::kernel::function_name>() == fname;
+  if (!res)
+    std::cout
+        << FFTestMark
+        << "test_kernel_bundle_ctxt failed: bundle does not contain kernel id "
+           "or function name "
+        << fname << std::endl;
+  return res;
+}
+
+template <auto *Func>
+int test_kernel_bundle_ctxt_dev(sycl::context &ctxt, sycl::device &dev,
+                                std::string_view fname) {
+  sycl::kernel_id id = syclexp::get_kernel_id<Func>();
+  auto exe_bndl =
+      syclexp::get_kernel_bundle<Func, sycl::bundle_state::executable>(ctxt,
+                                                                       {dev});
+  const bool res =
+      exe_bndl.has_kernel(id) &&
+      exe_bndl.get_kernel(id)
+              .template get_info<sycl::info::kernel::function_name>() == fname;
+  if (!res)
+    std::cout << FFTestMark
+              << "test_kernel_bundle_ctxt_dev failed: bundle does not contain "
+                 "kernel id "
+                 "or function name "
+              << fname << std::endl;
+  return res;
+}
+
+int main() {
+  sycl::queue q;
+  sycl::context ctxt = q.get_context();
+  sycl::device dev = q.get_device();
+
+  int ret = test_free_function_kernel_id<func_range>(ctxt);
+  ret |= test_free_function_kernel_id<func_single>(ctxt);
+  ret |= test_kernel_bundle_ctxt<func_range>(ctxt, "func_range");
+  ret |= test_kernel_bundle_ctxt<func_single>(ctxt, "func_single");
+  ret |= test_kernel_bundle_ctxt_dev<func_range>(ctxt, dev, "func_range");
+  ret |= test_kernel_bundle_ctxt_dev<func_single>(ctxt, dev, "func_single");
+  return ret;
+}

sycl/test-e2e/FreeFunctionKernels/separate_binaries.cpp

@@ -0,0 +1,38 @@
+// RUN: %{build} -fsycl-device-code-split=per_kernel -o %t.out
+// RUN: %{run} %t.out
+
+#include <iostream>
+#include <sycl/detail/core.hpp>
+#include <sycl/ext/oneapi/get_kernel_info.hpp>
+#include <sycl/kernel_bundle.hpp>
+
+namespace syclexp = sycl::ext::oneapi::experimental;
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<2>))
+void func_range(float start, float *ptr) {}
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::single_task_kernel))
+void func_single(float start, float *ptr) {}
+
+int main() {
+  sycl::queue q;
+  auto bundle_range =
+      syclexp::get_kernel_bundle<func_range, sycl::bundle_state::executable>(
+          q.get_context());
+  try {
+    sycl::kernel k = bundle_range.ext_oneapi_get_kernel<func_single>();
+  } catch (const sycl::exception &e) {
+    assert(e.code() == sycl::errc::invalid);
+    return 0;
+  }
+  auto bundle_single =
+      syclexp::get_kernel_bundle<func_single, sycl::bundle_state::executable>(
+          q.get_context());
+  try {
+    sycl::kernel k = bundle_single.ext_oneapi_get_kernel<func_range>();
+  } catch (const sycl::exception &e) {
+    assert(e.code() == sycl::errc::invalid);
+    return 0;
+  }
+  return 1;
+}