[SYCL][Graph] Node Profiling

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc

@@ -426,6 +426,13 @@ public:
  void ext_oneapi_graph(command_graph<graph_state::executable>& graph);
 }
 
+// New methods added to the sycl::event class
+class event {
+public:
+ template <typename Param>
+ typename Param::return_type ext_oneapi_get_profiling_info(node Node) const;
+}
+
 } // namespace sycl
 ----
 
@@ -1107,6 +1114,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
+ command groups associated with the event, then the call to this member
+ function will block until the requested info is available.
+
+Parameters:
+
+* `Node` - Handle to the node for which the profiling information is queried.
+
+Exceptions:
+
+* Throws synchronously with error code `invalid` if this SYCL event is not
+ associated to a graph execution.
+* Throws synchronously with error code `invalid` if the queue on which 
+ the graph is submitted was not constructed with 
+ the `property::queue::enable_profiling` property.
+* Throws synchronously with error code `invalid` if `Node` is not associated 
+ to the graph that was executed by event.
+
+|===
+
+
 === Thread Safety
 
 The new functions in this extension are thread-safe, the same as member
@@ -1810,8 +1855,10 @@ if used in application code.
 . Using `handler::memadvise` in a graph node.
 . 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::get_profiling_info(ext::node)` is only avalaible for
+ the level-zero backend.
 . 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

sycl/include/sycl/detail/pi.def

@@ -181,6 +181,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

@@ -2571,6 +2571,21 @@ 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 return 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 to 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 is the handle to the node for which the profiling information
+ /// is 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

@@ -1185,6 +1185,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

@@ -1193,6 +1193,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 param_name,
+ size_t param_value_size, void *param_value, size_t *param_value_size_ret) {
+ return pi2ur::piEventGetProfilingInfo(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/level_zero/pi_level_zero.cpp

@@ -1351,6 +1351,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

@@ -1176,6 +1176,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

@@ -1129,6 +1129,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,15 +56,9 @@ endif()
 if(SYCL_PI_UR_USE_FETCH_CONTENT)
  include(FetchContent)
 
- set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git")
- # commit 8d3233af35ed5515dd966c08d1af04b76113d3a8
- # Merge: eddda123 2bd41d7b
- # Author: Kenneth Benzie (Benie) <k.benzie@codeplay.com>
- # Date: Wed Jan 24 14:27:54 2024 +0000
- # Merge pull request #1247 from PietroGhg/pietro/empty_cmd_buffer_fill
- # [NATIVECPU] Empty implementation for command buffer fill
- set(UNIFIED_RUNTIME_TAG 8d3233af35ed5515dd966c08d1af04b76113d3a8)
-
+ set(UNIFIED_RUNTIME_REPO "https://github.com/bensuo/unified-runtime.git")
+ set(UNIFIED_RUNTIME_TAG cmdbuf-node-profiling)
+
  if(SYCL_PI_UR_OVERRIDE_FETCH_CONTENT_REPO)
  set(UNIFIED_RUNTIME_REPO "${SYCL_PI_UR_OVERRIDE_FETCH_CONTENT_REPO}")
  endif()

sycl/plugins/unified_runtime/pi2ur.hpp

@@ -4137,6 +4137,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(urSyncPointGetProfilingInfoExp(UREvent, SyncPoint, PropName,
+ ParamValueSize, ParamValue,
+ ParamValueSizeRet));
+
+ return PI_SUCCESS;
+}
+
 inline pi_result piEventCreate(pi_context Context, pi_event *RetEvent) {
 
  ur_context_handle_t UrContext =