Merge from 'sycl' to 'sycl-web'

Author: iclsrc

llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp

@@ -31,6 +31,7 @@ namespace id = itanium_demangle;
 #define DEBUG_TYPE "esimd-verifier"
 
 // A list of SYCL functions (regexps) allowed for use in ESIMD context.
+// clang-format off
 static const char *LegalSYCLFunctions[] = {
     "^sycl::_V1::accessor<.+>::accessor",
     "^sycl::_V1::accessor<.+>::~accessor",
@@ -80,6 +81,7 @@ static const char *LegalSYCLFunctionsInStatelessMode[] = {
     "^sycl::_V1::accessor<.+>::getMemoryRange",
     "^sycl::_V1::accessor<.+>::getOffset",
     "^sycl::_V1::accessor<.+>::operator\\[\\]"};
+// clang-format on
 
 namespace {
 

sycl/doc/design/DeviceLibExtensions.rst

@@ -4,7 +4,7 @@ Device library extensions
 Device compiler that indicates support for a particular extension is
 supposed to support *all* the corresponding functions.
 
-cl_intel_devicelib_cassert
+cl_intel_devicelib_assert
 ==========================
 
 .. code:

sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler_opencl.asciidoc

@@ -37,13 +37,13 @@ https://github.com/intel/llvm/issues
 
 == Dependencies
 
-This extension is written against the SYCL 2020 revision 7 specification.
+This extension is written against the SYCL 2020 revision 8 specification.
 All references below to the "core SYCL specification" or to section numbers in
 the SYCL specification refer to that revision.
 
 This extension also depends on the following other SYCL extensions:
 
-* link:../proposed/sycl_ext_oneapi_kernel_compiler.asciidoc[
+* link:../experimental/sycl_ext_oneapi_kernel_compiler.asciidoc[
   sycl_ext_oneapi_kernel_compiler]
 
 
@@ -60,7 +60,7 @@ this specification.*
 == Overview
 
 This is an extension to
-link:../proposed/sycl_ext_oneapi_kernel_compiler.asciidoc[
+link:../experimental/sycl_ext_oneapi_kernel_compiler.asciidoc[
 sycl_ext_oneapi_kernel_compiler], which allows an application to define a
 kernel in the OpenCL C language when dynamically compiling a kernel from
 source.

sycl/doc/extensions/proposed/sycl_ext_oneapi_free_function_queries.asciidoc

@@ -82,7 +82,7 @@ existing interfaces is not feasible.
 The `sycl::ext::oneapi::this_kernel` namespace contains functionality related
 to the currently executing kernel.
 
-It is the user's responsibility to ensure that that these functions are called
+It is the user's responsibility to ensure that these functions are called
 in a manner that is compatible with the kernel's launch parameters, as detailed
 in the definition of each function. Calling these functions from an incompatible
 kernel results in undefined behavior.

sycl/plugins/unified_runtime/CMakeLists.txt

@@ -57,13 +57,13 @@ if(SYCL_PI_UR_USE_FETCH_CONTENT)
   include(FetchContent)
 
   set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git")
-  # commit 71957e849435169b43e51a42025afeeb63d7508b
-  # Merge: ce152a64 e001b98d
+  # commit 31b654f981f6098936e7f04c65803395a2ea343a
+  # Merge: 71957e84 3da21336
   # Author: Kenneth Benzie (Benie) <k.benzie@codeplay.com>
-  # Date:   Tue Nov 21 16:41:08 2023 +0000
-  #     Merge pull request #1100 from kbenzie/benie/fixup-multi-device-compile
-  #     [CUDA][HIP][OpenCL][NATIVECPU] Fix multi-device compile
-  set(UNIFIED_RUNTIME_TAG 71957e849435169b43e51a42025afeeb63d7508b)
+  # Date:   Wed Nov 22 11:27:33 2023 +0000
+  #     Merge pull request #1053 from jandres742/url0leakkey
+  #     [UR][L0] Add UR_L0_LEAKS_DEBUG key
+  set(UNIFIED_RUNTIME_TAG 31b654f981f6098936e7f04c65803395a2ea343a)
 
   if(SYCL_PI_UR_OVERRIDE_FETCH_CONTENT_REPO)
     set(UNIFIED_RUNTIME_REPO "${SYCL_PI_UR_OVERRIDE_FETCH_CONTENT_REPO}")

sycl/source/detail/common.cpp

@@ -72,18 +72,24 @@ const char *stringifyErrorCode(pi_int32 error) {
 }
 
 std::vector<std::string> split_string(const std::string &str, char delimeter) {
-  std::vector<std::string> result;
-  size_t beg = 0;
-  size_t end = 0;
-  while ((end = str.find(delimeter, beg)) != std::string::npos) {
-    result.push_back(str.substr(beg, end - beg));
-    beg = end + 1;
+  std::vector<std::string> Result;
+  size_t Start = 0;
+  size_t End = 0;
+  while ((End = str.find(delimeter, Start)) != std::string::npos) {
+    Result.push_back(str.substr(Start, End - Start));
+    Start = End + 1;
   }
-  end = str.find('\0');
-  if (beg < end) {
-    result.push_back(str.substr(beg, end - beg));
+  // Get the last substring and ignore the null character so we wouldn't get
+  // double null characters \0\0 at the end of the substring
+  End = str.find('\0');
+  if (Start < End) {
+    std::string LastSubStr(str.substr(Start, End - Start));
+    // In case str has a delimeter at the end, the substring will be empty, so
+    // we shouldn't add it to the final vector
+    if (!LastSubStr.empty())
+      Result.push_back(LastSubStr);
   }
-  return result;
+  return Result;
 }
 
 } // namespace detail

sycl/source/detail/graph_impl.cpp

@@ -416,15 +416,13 @@ void graph_impl::makeEdge(std::shared_ptr<node_impl> Src,
 
 std::vector<sycl::detail::EventImplPtr> graph_impl::getExitNodesEvents() {
   std::vector<sycl::detail::EventImplPtr> Events;
-  auto EnqueueExitNodesEvents = [&](std::shared_ptr<node_impl> &Node,
-                                    std::deque<std::shared_ptr<node_impl>> &) {
+
+  for (auto Node : MNodeStorage) {
     if (Node->MSuccessors.empty()) {
       Events.push_back(getEventForNode(Node));
     }
-    return false;
-  };
+  }
 
-  searchDepthFirst(EnqueueExitNodesEvents);
   return Events;
 }
 

sycl/source/detail/graph_impl.hpp

@@ -784,6 +784,24 @@ class graph_impl {
   /// @return vector of events associated to exit nodes.
   std::vector<sycl::detail::EventImplPtr> getExitNodesEvents();
 
+  /// Removes all Barrier nodes from the list of extra dependencies
+  /// MExtraDependencies.
+  /// @return vector of events associated to previous barrier nodes.
+  std::vector<sycl::detail::EventImplPtr>
+  removeBarriersFromExtraDependencies() {
+    std::vector<sycl::detail::EventImplPtr> Events;
+    for (auto It = MExtraDependencies.begin();
+         It != MExtraDependencies.end();) {
+      if ((*It)->MCGType == sycl::detail::CG::Barrier) {
+        Events.push_back(getEventForNode(*It));
+        It = MExtraDependencies.erase(It);
+      } else {
+        ++It;
+      }
+    }
+    return Events;
+  }
+
 private:
   /// Iterate over the graph depth-first and run \p NodeFunc on each node.
   /// @param NodeFunc A function which receives as input a node in the graph to
@@ -861,7 +879,7 @@ class graph_impl {
   /// added to this graph.
   /// This list is mainly used by barrier nodes which must be considered
   /// as predecessors for all nodes subsequently added to the graph.
-  std::vector<std::shared_ptr<node_impl>> MExtraDependencies;
+  std::list<std::shared_ptr<node_impl>> MExtraDependencies;
 };
 
 /// Class representing the implementation of command_graph<executable>.

sycl/source/detail/queue_impl.cpp

@@ -85,6 +85,12 @@ event queue_impl::memset(const std::shared_ptr<detail::queue_impl> &Self,
   // Emit a begin/end scope for this call
   PrepareNotify.scopedNotify((uint16_t)xpti::trace_point_type_t::task_begin);
 #endif
+  if (MGraph.lock()) {
+    throw sycl::exception(make_error_code(errc::invalid),
+                          "The memset feature is not yet available "
+                          "for use with the SYCL Graph extension.");
+  }
+
   if (MHasDiscardEventsSupport) {
     MemoryManager::fill_usm(Ptr, Self, Count, Value,
                             getOrWaitEvents(DepEvents, MContext), nullptr);

sycl/source/handler.cpp

@@ -389,6 +389,14 @@ event handler::finalize() {
       // nodes/events of the graph
       if (MEventsWaitWithBarrier.size() == 0) {
         MEventsWaitWithBarrier = GraphImpl->getExitNodesEvents();
+        // Graph-wide barriers take precedence over previous one.
+        // We therefore remove the previous ones from ExtraDependencies list.
+        // The current barrier is then added to this list in the graph_impl.
+        std::vector<detail::EventImplPtr> EventsBarriers =
+            GraphImpl->removeBarriersFromExtraDependencies();
+        MEventsWaitWithBarrier.insert(std::end(MEventsWaitWithBarrier),
+                                      std::begin(EventsBarriers),
+                                      std::end(EventsBarriers));
       }
       CGData.MEvents.insert(std::end(CGData.MEvents),
                             std::begin(MEventsWaitWithBarrier),