[SYCL][Graph] Support for Graph Node Profiling

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc

@@ -1182,6 +1182,44 @@ Exceptions:
 
 |===
 
+==== New Event Member Functions
+
+Table {counter: tableNumber}. Additional member functions of the `sycl::event` class.
+[cols="2a,a"]
+|===
+|Member function|Description
+
+|
+[source,c++]
+----
+template <typename Param>
+typename Param::return_type 
+event::ext_oneapi_get_profiling_info(node Node) const;
+----
+
+| Queries the profiling information of a SYCL Graph node for the graph
+  execution associated with this SYCL event. If the requested info is
+  not available when this member function is called due to incompletion of
+  graph execution associated with the event, then the call to this member
+  function will block until the requested info is available.
+
+Parameters:
+
+* `Node` - Node object for which the profiling information is being queried.
+
+Exceptions:
+
+* Throws synchronously with error code `invalid` if this SYCL event is not
+  associated with a graph execution.
+* Throws synchronously with error code `invalid` if the queue on which 
+  the graph was submitted was not constructed with 
+  the `property::queue::enable_profiling` property.
+* Throws synchronously with error code `invalid` if `Node` is not associated 
+  with the graph exectution represented by this event.
+
+|===
+
+
 === Thread Safety
 
 The new functions in this extension are thread-safe, the same as member
@@ -1881,8 +1919,10 @@ if used in application code.
 
 . Using reductions in a graph node.
 . Using sycl streams in a graph node.
-. Profiling an event returned from graph submission with
-  `event::get_profiling_info()`.
+. Profiling information is not available for graphs that contain host-task nodes.
+. Profiling a node from an event returned from graph submission with
+  `event::ext_oneapi_get_profiling_info(ext::node)` is only available for
+  the level-zero backend.
 . Synchronization between multiple executions of the same command-buffer 
   must be handled in the host for level-zero backend, which may involve 
   extra latency for subsequent submissions.

sycl/include/sycl/detail/pi.def

@@ -183,6 +183,7 @@ _PI_API(piextCommandBufferFillUSM)
 _PI_API(piextCommandBufferPrefetchUSM)
 _PI_API(piextCommandBufferAdviseUSM)
 _PI_API(piextEnqueueCommandBuffer)
+_PI_API(piextSyncPointGetProfilingInfo)
 
 _PI_API(piextUSMPitchedAlloc)
 

sycl/include/sycl/detail/pi.h

@@ -154,9 +154,10 @@
 // 15.44 Add coarse-grain memory advice flag for HIP.
 // 15.45 Added piextKernelSuggestMaxCooperativeGroupCount and
 //       piextEnqueueCooperativeKernelLaunch.
+// 15.46 Added piextSyncPointGetProfilingInfo
 
 #define _PI_H_VERSION_MAJOR 15
-#define _PI_H_VERSION_MINOR 45
+#define _PI_H_VERSION_MINOR 46
 
 #define _PI_STRING_HELPER(a) #a
 #define _PI_CONCAT(a, b) _PI_STRING_HELPER(a.b)
@@ -2601,6 +2602,24 @@ piextEnqueueCommandBuffer(pi_ext_command_buffer command_buffer, pi_queue queue,
                           pi_uint32 num_events_in_wait_list,
                           const pi_event *event_wait_list, pi_event *event);
 
+/// API to get the profiling information of a graph node.
+/// A Node is identified by a sync-point in a command-buffer.
+/// The sync-point passed in parameter corresponds therefore to the node from
+/// which we want to get the profiling information. returns an error if the node
+/// is found.
+/// \param event PI event that has been returned from the command-buffer
+/// submission.
+/// \param sync_point The sync-point corresponding to the node from which
+/// we want to get the profiling information.
+/// \param param_name The name of the profiling property to query depends on.
+/// \param param_value_size Size in bytes of the profiling property value.
+/// \param param_value Value of the profiling property.
+/// \param param_value_size_ret Pointer to the actual size in bytes returned
+/// in param_value of the profiling property.
+__SYCL_EXPORT pi_result piextSyncPointGetProfilingInfo(
+    pi_event event, pi_ext_sync_point sync_point, pi_profiling_info param_name,
+    size_t param_value_size, void *param_value, size_t *param_value_size_ret);
+
 /// API to destroy bindless unsampled image handles.
 ///
 /// \param context is the pi_context

sycl/include/sycl/event.hpp

@@ -14,6 +14,7 @@
 #include <sycl/detail/info_desc_helpers.hpp>  // for is_event_info_desc, is_...
 #include <sycl/detail/owner_less_base.hpp>    // for OwnerLessBase
 #include <sycl/detail/pi.h>                   // for pi_native_handle
+#include <sycl/ext/oneapi/experimental/graph.hpp>
 
 #ifdef __SYCL_INTERNAL_API
 #include <sycl/detail/cl.h>
@@ -130,6 +131,25 @@ class __SYCL_EXPORT event : public detail::OwnerLessBase<event> {
   typename detail::is_event_profiling_info_desc<Param>::return_type
   get_profiling_info() const;
 
+  /// Queries the profiling information of a SYCL Graph node for the graph
+  /// execution associated with this SYCL event.
+  ///
+  /// If this SYCL event is not associated with a graph execution, an
+  /// invalid_object_error SYCL exception is thrown. If the requested info is
+  /// not available when this member function is called due to incompletion of
+  /// command groups associated with the event, then the call to this member
+  /// function will block until the requested info is available. If the queue
+  /// which submitted the command group this event is associated with was not
+  /// constructed with the property::queue::enable_profiling property, an
+  /// invalid_object_error SYCL exception is thrown.
+  ///
+  /// \param Node Node object for which the profiling information
+  /// is being queried.
+  /// \return depends on template parameter.
+  template <typename Param>
+  typename detail::is_event_profiling_info_desc<Param>::return_type
+  ext_oneapi_get_profiling_info(ext::oneapi::experimental::node Node) const;
+
   /// Returns the backend associated with this platform.
   ///
   /// \return the backend associated with this platform

sycl/include/sycl/info/ext_oneapi_graph_node_profiling_traits.def

@@ -0,0 +1,3 @@
+__SYCL_PARAM_TRAITS_SPEC_PARAMT(event_profiling, command_submit, ext::oneapi::experimental::node, uint64_t, PI_PROFILING_INFO_COMMAND_SUBMIT)
+__SYCL_PARAM_TRAITS_SPEC_PARAMT(event_profiling, command_start, ext::oneapi::experimental::node, uint64_t, PI_PROFILING_INFO_COMMAND_START)
+__SYCL_PARAM_TRAITS_SPEC_PARAMT(event_profiling, command_end, ext::oneapi::experimental::node, uint64_t, PI_PROFILING_INFO_COMMAND_END)

sycl/include/sycl/info/info_desc.hpp

@@ -36,6 +36,11 @@ namespace info {
   struct Desc {                                                                \
     using return_type = ReturnT;                                               \
   };
+#define __SYCL_PARAM_TRAITS_SPEC_PARAMT(DescType, Desc, ParamType, ReturnT,    \
+                                        PiCode)                                \
+  struct Desc {                                                                \
+    using return_type = ReturnT;                                               \
+  };
 // A.1 Platform information desctiptors
 namespace platform {
 // TODO Despite giving this deprecation warning, we're still yet to implement
@@ -155,7 +160,11 @@ namespace event {
 namespace event_profiling {
 #include <sycl/info/event_profiling_traits.def>
 } // namespace event_profiling
+namespace ext_oneapi_event_profiling {
+#include <sycl/info/ext_oneapi_graph_node_profiling_traits.def>
+} // namespace ext_oneapi_event_profiling
 #undef __SYCL_PARAM_TRAITS_SPEC
+#undef __SYCL_PARAM_TRAITS_SPEC_PARAMT
 
 // Provide an alias to the return type for each of the info parameters
 template <typename T, T param> class param_traits {};

sycl/plugins/cuda/pi_cuda.cpp

@@ -1208,6 +1208,14 @@ pi_result piextEnqueueCommandBuffer(pi_ext_command_buffer CommandBuffer,
       CommandBuffer, Queue, NumEventsInWaitList, EventWaitList, Event);
 }
 
+pi_result piextSyncPointGetProfilingInfo(
+    pi_event Event, pi_ext_sync_point SyncPoint, pi_profiling_info ParamName,
+    size_t ParamValueSize, void *ParamValue, size_t *ParamValueSizeRet) {
+  return pi2ur::piextSyncPointGetProfilingInfo(Event, SyncPoint, ParamName,
+                                               ParamValueSize, ParamValue,
+                                               ParamValueSizeRet);
+}
+
 pi_result piextPluginGetOpaqueData(void *opaque_data_param,
                                    void **opaque_data_return) {
   return pi2ur::piextPluginGetOpaqueData(opaque_data_param, opaque_data_return);

sycl/plugins/hip/pi_hip.cpp

@@ -1211,6 +1211,14 @@ pi_result piextEnqueueCommandBuffer(pi_ext_command_buffer CommandBuffer,
       CommandBuffer, Queue, NumEventsInWaitList, EventWaitList, Event);
 }
 
+pi_result piextSyncPointGetProfilingInfo(
+    pi_event event, pi_ext_sync_point sync_point, pi_profiling_info param_name,
+    size_t param_value_size, void *param_value, size_t *param_value_size_ret) {
+  return pi2ur::piextSyncPointGetProfilingInfo(event, sync_point, param_name,
+                                               param_value_size, param_value,
+                                               param_value_size_ret);
+}
+
 pi_result piextPluginGetOpaqueData(void *opaque_data_param,
                                    void **opaque_data_return) {
   return pi2ur::piextPluginGetOpaqueData(opaque_data_param, opaque_data_return);

sycl/plugins/level_zero/pi_level_zero.cpp

@@ -1369,6 +1369,14 @@ pi_result piextEnqueueCommandBuffer(pi_ext_command_buffer CommandBuffer,
       CommandBuffer, Queue, NumEventsInWaitList, EventWaitList, Event);
 }
 
+pi_result piextSyncPointGetProfilingInfo(
+    pi_event Event, pi_ext_sync_point SyncPoint, pi_profiling_info ParamName,
+    size_t ParamValueSize, void *ParamValue, size_t *ParamValueSizeRet) {
+  return pi2ur::piextSyncPointGetProfilingInfo(Event, SyncPoint, ParamName,
+                                               ParamValueSize, ParamValue,
+                                               ParamValueSizeRet);
+}
+
 const char SupportedVersion[] = _PI_LEVEL_ZERO_PLUGIN_VERSION_STRING;
 
 pi_result piPluginInit(pi_plugin *PluginInit) { // missing

sycl/plugins/native_cpu/pi_native_cpu.cpp

@@ -1177,6 +1177,14 @@ pi_result piextEnqueueCommandBuffer(pi_ext_command_buffer CommandBuffer,
       CommandBuffer, Queue, NumEventsInWaitList, EventWaitList, Event);
 }
 
+pi_result piextSyncPointGetProfilingInfo(
+    pi_event Event, pi_ext_sync_point SyncPoint, pi_profiling_info ParamName,
+    size_t ParamValueSize, void *ParamValue, size_t *ParamValueSizeRet) {
+  return pi2ur::piextSyncPointGetProfilingInfo(Event, SyncPoint, ParamName,
+                                               ParamValueSize, ParamValue,
+                                               ParamValueSizeRet);
+}
+
 pi_result piextPluginGetOpaqueData(void *opaque_data_param,
                                    void **opaque_data_return) {
   return pi2ur::piextPluginGetOpaqueData(opaque_data_param, opaque_data_return);

sycl/plugins/opencl/pi_opencl.cpp

@@ -1147,6 +1147,14 @@ pi_result piextEnqueueCommandBuffer(pi_ext_command_buffer CommandBuffer,
       CommandBuffer, Queue, NumEventsInWaitList, EventWaitList, Event);
 }
 
+pi_result piextSyncPointGetProfilingInfo(
+    pi_event Event, pi_ext_sync_point SyncPoint, pi_profiling_info ParamName,
+    size_t ParamValueSize, void *ParamValue, size_t *ParamValueSizeRet) {
+  return pi2ur::piextSyncPointGetProfilingInfo(Event, SyncPoint, ParamName,
+                                               ParamValueSize, ParamValue,
+                                               ParamValueSizeRet);
+}
+
 pi_result piextPluginGetOpaqueData(void *opaque_data_param,
                                    void **opaque_data_return) {
   return pi2ur::piextPluginGetOpaqueData(opaque_data_param, opaque_data_return);

sycl/plugins/unified_runtime/CMakeLists.txt

@@ -56,14 +56,8 @@ endif()
 if(SYCL_PI_UR_USE_FETCH_CONTENT)
   include(FetchContent)
 
-  set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git")
-  # commit e2ee9a4720414e0a59fa9c911e9575ab564ac57c
-  # Merge: 7a5150cd be622e7c
-  # Author: Kenneth Benzie (Benie) <k.benzie@codeplay.com>
-  # Date:   Sun Mar 10 18:02:50 2024 +0000
-  #     Merge pull request #1340 from Bensuo/ewan/coverity_cuda_update
-  #     [HIP][CUDA][Command-Buffer] Fix Coverity issues in HIP/CUDA command-buffer code
-  set(UNIFIED_RUNTIME_TAG e2ee9a4720414e0a59fa9c911e9575ab564ac57c)
+  set(UNIFIED_RUNTIME_REPO "https://github.com/bensuo/unified-runtime.git")
+  set(UNIFIED_RUNTIME_TAG cmdbuf-profiling-sync-point)
 
   if(SYCL_PI_UR_OVERRIDE_FETCH_CONTENT_REPO)
     set(UNIFIED_RUNTIME_REPO "${SYCL_PI_UR_OVERRIDE_FETCH_CONTENT_REPO}")

sycl/plugins/unified_runtime/pi2ur.hpp

@@ -4235,6 +4235,43 @@ inline pi_result piEventGetProfilingInfo(pi_event Event,
   return PI_SUCCESS;
 }
 
+inline pi_result piextSyncPointGetProfilingInfo(
+    pi_event Event, pi_ext_sync_point SyncPoint, pi_profiling_info ParamName,
+    size_t ParamValueSize, void *ParamValue, size_t *ParamValueSizeRet) {
+
+  PI_ASSERT(Event, PI_ERROR_INVALID_EVENT);
+
+  ur_event_handle_t UREvent = reinterpret_cast<ur_event_handle_t>(Event);
+
+  ur_profiling_info_t PropName{};
+  switch (ParamName) {
+  case PI_PROFILING_INFO_COMMAND_QUEUED: {
+    PropName = UR_PROFILING_INFO_COMMAND_QUEUED;
+    break;
+  }
+  case PI_PROFILING_INFO_COMMAND_SUBMIT: {
+    PropName = UR_PROFILING_INFO_COMMAND_SUBMIT;
+    break;
+  }
+  case PI_PROFILING_INFO_COMMAND_START: {
+    PropName = UR_PROFILING_INFO_COMMAND_START;
+    break;
+  }
+  case PI_PROFILING_INFO_COMMAND_END: {
+    PropName = UR_PROFILING_INFO_COMMAND_END;
+    break;
+  }
+  default:
+    return PI_ERROR_INVALID_PROPERTY;
+  }
+
+  HANDLE_ERRORS(urEventGetSyncPointProfilingInfoExp(
+      UREvent, SyncPoint, PropName, ParamValueSize, ParamValue,
+      ParamValueSizeRet));
+
+  return PI_SUCCESS;
+}
+
 inline pi_result piEventCreate(pi_context Context, pi_event *RetEvent) {
 
   ur_context_handle_t UrContext =