Merge from 'sycl' to 'sycl-web'

Author: iclsrc

libclc/amdgcn-amdhsa/libspirv/atomic/atomic_cmpxchg.cl

@@ -23,12 +23,17 @@
                                memory_order_success)                                                                                                                 \
     GET_ATOMIC_SCOPE_AND_ORDER(scope, atomic_scope, failure_semantics,                                                                                               \
                                memory_order_failure)                                                                                                                 \
-    TYPE original_val = *p;                                                                                                                                        \
-    bool success = __hip_atomic_compare_exchange_strong(                                                                                                             \
-        p, &expected, desired, memory_order_success, memory_order_failure,                                                                                           \
-        atomic_scope);                                                                                                                                               \
-                                                                                                                                                                     \
-    return success ? original_val : *p;                                                                                                                              \
+    __hip_atomic_compare_exchange_strong(p, &expected, desired,                                                                                                      \
+                                         memory_order_success,                                                                                                       \
+                                         memory_order_failure, atomic_scope);                                                                                        \
+    /* If cmpxchg                                                                                                                                                    \
+     *  succeeds:                                                                                                                                                    \
+         - `expected` is unchanged, holding the old val that was at `p`                                                                                              \
+         - `p` is changed to hold `desired`                                                                                                                          \
+     *  fails:                                                                                                                                                       \
+         - `expected` is changed to hold the current val at `p`                                                                                                      \
+         - `p` is unchanged*/                                                                                                                                        \
+    return expected;                                                                                                                                                 \
   }
 
 #define AMDGPU_ATOMIC_CMPXCHG(TYPE, TYPE_MANGLED)                              \
@@ -37,7 +42,7 @@
   AMDGPU_ATOMIC_CMPXCHG_IMPL(TYPE, TYPE_MANGLED, , , 0, 4)
 
 AMDGPU_ATOMIC_CMPXCHG(int, i)
-AMDGPU_ATOMIC_CMPXCHG(unsigned int, j)
+AMDGPU_ATOMIC_CMPXCHG(unsigned, j)
 AMDGPU_ATOMIC_CMPXCHG(long, l)
 AMDGPU_ATOMIC_CMPXCHG(unsigned long, m)
 AMDGPU_ATOMIC_CMPXCHG(float, f)

libclc/amdgcn-amdhsa/libspirv/atomic/atomic_helpers.h

@@ -71,3 +71,25 @@
   AMDGPU_ATOMIC_IMPL(FUNC_NAME, TYPE, TYPE_MANGLED, global, U3AS1, 1, BUILTIN) \
   AMDGPU_ATOMIC_IMPL(FUNC_NAME, TYPE, TYPE_MANGLED, local, U3AS3, 1, BUILTIN)  \
   AMDGPU_ATOMIC_IMPL(FUNC_NAME, TYPE, TYPE_MANGLED, , , 0, BUILTIN)
+
+#define AMDGPU_CAS_ATOMIC_IMPL(FUNC_NAME, TYPE, TYPE_MANGLED, AS, AS_MANGLED,                                          \
+                               SUB1, OP)                                                                               \
+  _CLC_DEF TYPE                                                                                                        \
+      FUNC_NAME##P##AS_MANGLED##TYPE_MANGLED##N5__spv5Scope4FlagENS##SUB1##_19MemorySemanticsMask4FlagE##TYPE_MANGLED( \
+          volatile AS TYPE *p, enum Scope scope,                                                                       \
+          enum MemorySemanticsMask semantics, TYPE val) {                                                              \
+    int atomic_scope = 0, memory_order = 0;                                                                            \
+    GET_ATOMIC_SCOPE_AND_ORDER(scope, atomic_scope, semantics, memory_order)                                           \
+    TYPE oldval = __hip_atomic_load(p, memory_order, atomic_scope);                                                    \
+    TYPE newval = 0;                                                                                                   \
+    do {                                                                                                               \
+      newval = oldval OP val;                                                                                          \
+    } while (!__hip_atomic_compare_exchange_strong(                                                                    \
+        p, &oldval, newval, atomic_scope, atomic_scope, memory_order));                                                \
+    return oldval;                                                                                                     \
+  }
+
+#define AMDGPU_CAS_ATOMIC(FUNC_NAME, TYPE, TYPE_MANGLED, OP)                   \
+  AMDGPU_CAS_ATOMIC_IMPL(FUNC_NAME, TYPE, TYPE_MANGLED, global, U3AS1, 1, OP)  \
+  AMDGPU_CAS_ATOMIC_IMPL(FUNC_NAME, TYPE, TYPE_MANGLED, local, U3AS3, 1, OP)   \
+  AMDGPU_CAS_ATOMIC_IMPL(FUNC_NAME, TYPE, TYPE_MANGLED, , , 0, OP)

libclc/amdgcn-amdhsa/libspirv/atomic/atomic_xor.cl

@@ -10,11 +10,14 @@
 #include <spirv/spirv.h>
 #include <spirv/spirv_types.h>
 
-AMDGPU_ATOMIC(_Z17__spirv_AtomicXor, int, i, __hip_atomic_fetch_xor)
-AMDGPU_ATOMIC(_Z17__spirv_AtomicXor, unsigned int, j, __hip_atomic_fetch_xor)
-AMDGPU_ATOMIC(_Z17__spirv_AtomicXor, long, l, __hip_atomic_fetch_xor)
-AMDGPU_ATOMIC(_Z17__spirv_AtomicXor, unsigned long, m, __hip_atomic_fetch_xor)
+#define __CLC_XOR ^
 
+AMDGPU_CAS_ATOMIC(_Z17__spirv_AtomicXor, int, i, __CLC_XOR)
+AMDGPU_CAS_ATOMIC(_Z17__spirv_AtomicXor, unsigned int, j, __CLC_XOR)
+AMDGPU_CAS_ATOMIC(_Z17__spirv_AtomicXor, long, l, __CLC_XOR)
+AMDGPU_CAS_ATOMIC(_Z17__spirv_AtomicXor, unsigned long, m, __CLC_XOR)
+
+#undef __CLC_XOR
 #undef AMDGPU_ATOMIC
 #undef AMDGPU_ATOMIC_IMPL
 #undef AMDGPU_ARCH_GEQ

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc

@@ -336,6 +336,11 @@ class depends_on {
     depends_on(NodeTN... nodes);
 };
 
+class depends_on_all_leaves {
+  public:
+    depends_on_all_leaves() = default;
+};
+
 } // namespace node
 } // namespace property
 
@@ -481,6 +486,21 @@ class depends_on {
 }
 ----
 
+==== Depends-On-All-Leaves Property
+The API for explicitly adding nodes to a `command_graph` includes a
+`property_list` parameter. This extension defines the `depends_on_all_leaves`
+property to be passed here. `depends_on_all_leaves` provides a shortcut for
+adding all the current leaves of a graph as dependencies.
+[source,c++]
+----
+namespace sycl::ext::oneapi::experimental::property::node {
+class depends_on_all_leaves {
+  public:
+    depends_on_all_leaves();
+};
+}
+----
+
 === Graph
 
 This extension adds a new `command_graph` object which follows the
@@ -737,11 +757,11 @@ print_graph(std::string path, bool verbose = false) const;
 ----
 
 |Synchronous operation that writes a DOT formatted description of the graph to the
-provided path. By default, this includes the graph topology, node types, node id, 
-and kernel names. 
-Verbose can be set to true to write more detailed information about each node type 
+provided path. By default, this includes the graph topology, node types, node id,
+and kernel names.
+Verbose can be set to true to write more detailed information about each node type
 such as kernel arguments, copy source, and destination addresses.
-At the moment DOT format is the only supported format. The name of hte output file 
+At the moment DOT format is the only supported format. The name of the output file
 must therefore match this extension, i.e. "<filename>.dot".
 
 Parameters:
@@ -752,7 +772,7 @@ or memory access where applicable.
 
 Exceptions:
 
-* Throws synchronously with error code `invalid` if the path is invalid or 
+* Throws synchronously with error code `invalid` if the path is invalid or
 the file extension is not supported or if the write operation failed.
 
 |===
@@ -1200,11 +1220,12 @@ passed an invalid event.
 
 The new handler methods, and queue shortcuts, defined by
 link:../supported/sycl_ext_oneapi_enqueue_barrier.asciidoc[sycl_ext_oneapi_enqueue_barrier]
-cannot be used in graph nodes. A synchronous exception will be thrown with
-error code `invalid` if a user tries to add them to a graph.
-
-Removing this restriction is something we may look at for future revisions of
-`sycl_ext_oneapi_graph`.
+can only be used in graph nodes created using the Record & Replay API, as
+barriers rely on events to enforce dependencies. A synchronous exception will be
+thrown with error code `invalid` if a user tries to add them to a graph using
+the Explicit API. Empty nodes created with the `node::depends_on_all_leaves`
+property can be used instead of barriers when a user is building a graph with
+the explicit API.
 
 ==== sycl_ext_oneapi_memcpy2d
 
@@ -1773,12 +1794,14 @@ if used in application code.
 . Using `handler::memset` in a graph node.
 . Using `handler::prefetch` in a graph node.
 . Using `handler::memadvise` in a graph node.
-. Using specialization constants in a graph node.
 . Using reductions in a graph node.
 . Using sycl streams in a graph node.
-. Using a kernel bundle in a graph node.
 . Profiling an event returned from graph submission with
   `event::get_profiling_info()`.
+. Level Zero immediate command-lists are not supported, and
+  `sycl::ext::intel::property::queue::no_immediate_command_list`
+  should be set on construction to any queues an executable
+  graph is submitted to.
 
 == Revision History
 

sycl/include/sycl/detail/property_helper.hpp

@@ -46,8 +46,9 @@ enum DataLessPropKind {
   QueueSubmissionImmediate = 21,
   GraphAssumeDataOutlivesBuffer = 22,
   GraphAssumeBufferOutlivesGraph = 23,
+  GraphDependOnAllLeaves = 24,
   // Indicates the last known dataless property.
-  LastKnownDataLessPropKind = 23,
+  LastKnownDataLessPropKind = 24,
   // Exceeding 32 may cause ABI breaking change on some of OSes.
   DataLessPropKindSize = 32
 };

sycl/include/sycl/ext/oneapi/experimental/graph.hpp

@@ -137,6 +137,14 @@ class depends_on : public ::sycl::detail::PropertyWithData<
   const std::vector<::sycl::ext::oneapi::experimental::node> MDeps;
 };
 
+/// Property used to to add all previous graph leaves as dependencies when
+/// creating a new node with command_graph::add().
+class depends_on_all_leaves : public ::sycl::detail::DataLessProperty<
+                                  ::sycl::detail::GraphDependOnAllLeaves> {
+public:
+  depends_on_all_leaves() = default;
+};
+
 } // namespace node
 } // namespace property
 
@@ -159,9 +167,17 @@ class __SYCL_EXPORT modifiable_command_graph {
   node add(const property_list &PropList = {}) {
     if (PropList.has_property<property::node::depends_on>()) {
       auto Deps = PropList.get_property<property::node::depends_on>();
-      return addImpl(Deps.get_dependencies());
+      node Node = addImpl(Deps.get_dependencies());
+      if (PropList.has_property<property::node::depends_on_all_leaves>()) {
+        addGraphLeafDependencies(Node);
+      }
+      return Node;
     }
-    return addImpl({});
+    node Node = addImpl({});
+    if (PropList.has_property<property::node::depends_on_all_leaves>()) {
+      addGraphLeafDependencies(Node);
+    }
+    return Node;
   }
 
   /// Add a command-group node to the graph.
@@ -171,9 +187,17 @@ class __SYCL_EXPORT modifiable_command_graph {
   template <typename T> node add(T CGF, const property_list &PropList = {}) {
     if (PropList.has_property<property::node::depends_on>()) {
       auto Deps = PropList.get_property<property::node::depends_on>();
-      return addImpl(CGF, Deps.get_dependencies());
+      node Node = addImpl(CGF, Deps.get_dependencies());
+      if (PropList.has_property<property::node::depends_on_all_leaves>()) {
+        addGraphLeafDependencies(Node);
+      }
+      return Node;
+    }
+    node Node = addImpl(CGF, {});
+    if (PropList.has_property<property::node::depends_on_all_leaves>()) {
+      addGraphLeafDependencies(Node);
     }
-    return addImpl(CGF, {});
+    return Node;
   }
 
   /// Add a dependency between two nodes.
@@ -247,6 +271,11 @@ class __SYCL_EXPORT modifiable_command_graph {
   /// @return Node added to the graph.
   node addImpl(const std::vector<node> &Dep);
 
+  /// Adds all graph leaves as dependencies
+  /// @param Node Destination node to which the leaves of the graph will be
+  /// added as dependencies.
+  void addGraphLeafDependencies(node Node);
+
   template <class Obj>
   friend decltype(Obj::impl)
   sycl::detail::getSyclObjImpl(const Obj &SyclObject);

sycl/plugins/cuda/CMakeLists.txt

@@ -63,5 +63,10 @@ if (SYCL_ENABLE_XPTI_TRACING)
   )
 endif()
 
+if(CUDA_cupti_LIBRARY)
+  target_compile_definitions(pi_cuda PRIVATE
+    "-DCUPTI_LIB_PATH=\"${CUDA_cupti_LIBRARY}\"")
+endif()
+
 set_target_properties(pi_cuda PROPERTIES LINKER_LANGUAGE CXX)
 

sycl/source/detail/graph_impl.cpp

@@ -726,6 +726,19 @@ node modifiable_command_graph::addImpl(std::function<void(handler &)> CGF,
   return sycl::detail::createSyclObjFromImpl<node>(NodeImpl);
 }
 
+void modifiable_command_graph::addGraphLeafDependencies(node Node) {
+  // Find all exit nodes in the current graph and add them to the dependency
+  // vector
+  std::shared_ptr<detail::node_impl> DstImpl =
+      sycl::detail::getSyclObjImpl(Node);
+  graph_impl::WriteLock Lock(impl->MMutex);
+  for (auto &NodeImpl : impl->MNodeStorage) {
+    if ((NodeImpl->MSuccessors.size() == 0) && (NodeImpl != DstImpl)) {
+      impl->makeEdge(NodeImpl, DstImpl);
+    }
+  }
+}
+
 void modifiable_command_graph::make_edge(node &Src, node &Dest) {
   std::shared_ptr<detail::node_impl> SenderImpl =
       sycl::detail::getSyclObjImpl(Src);

sycl/test/abi/sycl_symbols_linux.dump

@@ -3734,6 +3734,7 @@ _ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graph7addImplES
 _ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graph9make_edgeERNS3_4nodeES7_
 _ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graphC1ERKNS0_7contextERKNS0_6deviceERKNS0_13property_listE
 _ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graphC2ERKNS0_7contextERKNS0_6deviceERKNS0_13property_listE
+_ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graph24addGraphLeafDependenciesENS3_4nodeE
 _ZNK4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graph11print_graphENSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEEb
 _ZN4sycl3_V13ext6oneapi12experimental9image_memC1ERKNS3_16image_descriptorERKNS0_5queueE
 _ZN4sycl3_V13ext6oneapi12experimental9image_memC1ERKNS3_16image_descriptorERKNS0_6deviceERKNS0_7contextE

sycl/test/abi/sycl_symbols_windows.dump

@@ -873,6 +873,7 @@
 ?add@device_global_map@detail@_V1@sycl@@YAXPEBXPEBD@Z
 ?add@host_pipe_map@detail@_V1@sycl@@YAXPEBXPEBD@Z
 ?add@modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@QEAA?AVnode@34567@AEBVproperty_list@67@@Z
+?addGraphLeafDependencies@modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@IEAAXVnode@34567@@Z
 ?addHostAccessorAndWait@detail@_V1@sycl@@YAXPEAVAccessorImplHost@123@@Z
 ?addHostSampledImageAccessorAndWait@detail@_V1@sycl@@YAXPEAVSampledImageAccessorImplHost@123@@Z
 ?addHostUnsampledImageAccessorAndWait@detail@_V1@sycl@@YAXPEAVUnsampledImageAccessorImplHost@123@@Z