[SYCL] Implement sycl_ext_oneapi_profiling_tag extension (intel#12838)

This commit adds the implementation of the sycl_ext_oneapi_profiling_tag
extension. Moving the extension to experimental will happen in a
follow-up patch.

---------

Signed-off-by: Larsen, Steffen <[email protected]>

Author: steffenlarsen

llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td

@@ -81,6 +81,7 @@ def AspectExt_oneapi_graph : Aspect<"ext_oneapi_graph">;
 def AspectExt_intel_fpga_task_sequence : Aspect<"ext_intel_fpga_task_sequence">;
 def AspectExt_oneapi_limited_graph : Aspect<"ext_oneapi_limited_graph">;
 def AspectExt_oneapi_private_alloca : Aspect<"ext_oneapi_private_alloca">;
+def AspectExt_oneapi_queue_profiling_tag : Aspect<"ext_oneapi_queue_profiling_tag">;
 // Deprecated aspects
 def AspectInt64_base_atomics : Aspect<"int64_base_atomics">;
 def AspectInt64_extended_atomics : Aspect<"int64_extended_atomics">;
@@ -138,8 +139,8 @@ def : TargetInfo<"__TestAspectList",
     AspectExt_oneapi_ballot_group, AspectExt_oneapi_fixed_size_group, AspectExt_oneapi_opportunistic_group,
     AspectExt_oneapi_tangle_group, AspectExt_intel_matrix, AspectExt_oneapi_is_composite, AspectExt_oneapi_is_component,
     AspectExt_oneapi_graph, AspectExt_intel_fpga_task_sequence, AspectExt_oneapi_limited_graph,
-    AspectExt_oneapi_private_alloca],
-[]>;
+    AspectExt_oneapi_private_alloca, AspectExt_oneapi_queue_profiling_tag],
+    []>;
 // This definition serves the only purpose of testing whether the deprecated aspect list defined in here and in SYCL RT
 // match.
 def : TargetInfo<"__TestDeprecatedAspectList",

sycl/include/sycl/detail/cg.hpp

@@ -78,6 +78,7 @@ class CG {
     CopyImage = 23,
     SemaphoreWait = 24,
     SemaphoreSignal = 25,
+    ProfilingTag = 26,
   };
 
   struct StorageInitHelper {
@@ -344,6 +345,12 @@ class CGBarrier : public CG {
         MEventsWaitWithBarrier(std::move(EventsWaitWithBarrier)) {}
 };
 
+class CGProfilingTag : public CG {
+public:
+  CGProfilingTag(CG::StorageInitHelper CGData, detail::code_location loc = {})
+      : CG(CG::ProfilingTag, std::move(CGData), std::move(loc)) {}
+};
+
 /// "Copy 2D USM" command group class.
 class CGCopy2DUSM : public CG {
   void *MSrc;

sycl/include/sycl/detail/pi.def

@@ -100,6 +100,7 @@ _PI_API(piEventRetain)
 _PI_API(piEventRelease)
 _PI_API(piextEventGetNativeHandle)
 _PI_API(piextEventCreateWithNativeHandle)
+_PI_API(piEnqueueTimestampRecordingExp)
 // Sampler
 _PI_API(piSamplerCreate)
 _PI_API(piSamplerGetInfo)

sycl/include/sycl/detail/pi.h

@@ -174,9 +174,11 @@
 //         - PI_EXT_ONEAPI_DEVICE_INFO_BINDLESS_SAMPLED_IMAGE_FETCH_3D
 // 15.51 Removed ret_mem argument from piextMemUnsampledImageCreate and
 // piextMemSampledImageCreate
+// 15.52 Added piEnqueueTimestampRecordingExp and
+//       PI_EXT_ONEAPI_DEVICE_INFO_TIMESTAMP_RECORDING_SUPPORT.
 
 #define _PI_H_VERSION_MAJOR 15
-#define _PI_H_VERSION_MINOR 51
+#define _PI_H_VERSION_MINOR 52
 
 #define _PI_STRING_HELPER(a) #a
 #define _PI_CONCAT(a, b) _PI_STRING_HELPER(a.b)
@@ -484,6 +486,9 @@ typedef enum {
   PI_EXT_ONEAPI_DEVICE_INFO_BINDLESS_SAMPLED_IMAGE_FETCH_2D = 0x2011A,
   PI_EXT_ONEAPI_DEVICE_INFO_BINDLESS_SAMPLED_IMAGE_FETCH_3D_USM = 0x2011B,
   PI_EXT_ONEAPI_DEVICE_INFO_BINDLESS_SAMPLED_IMAGE_FETCH_3D = 0x2011C,
+
+  // Timestamp enqueue
+  PI_EXT_ONEAPI_DEVICE_INFO_TIMESTAMP_RECORDING_SUPPORT = 0x2011D,
 } _pi_device_info;
 
 typedef enum {
@@ -1774,6 +1779,10 @@ __SYCL_EXPORT pi_result piEventRetain(pi_event event);
 
 __SYCL_EXPORT pi_result piEventRelease(pi_event event);
 
+__SYCL_EXPORT pi_result piEnqueueTimestampRecordingExp(
+    pi_queue queue, pi_bool blocking, pi_uint32 num_events_in_wait_list,
+    const pi_event *event_wait_list, pi_event *event);
+
 /// Gets the native handle of a PI event object.
 ///
 /// \param event is the PI event to get the native handle of.

sycl/include/sycl/device_aspect_macros.hpp

@@ -376,6 +376,11 @@
 #define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_bindless_sampled_image_fetch_3d__ 0
 #endif
 
+#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_queue_profiling_tag__
+// __SYCL_ASPECT(ext_oneapi_queue_profiling_tag, 73)
+#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_queue_profiling_tag__ 0
+#endif
+
 #ifndef __SYCL_ANY_DEVICE_HAS_host__
 // __SYCL_ASPECT(host, 0)
 #define __SYCL_ANY_DEVICE_HAS_host__ 0
@@ -740,3 +745,8 @@
 //__SYCL_ASPECT(ext_oneapi_bindless_sampled_image_fetch_3d, 72)
 #define __SYCL_ANY_DEVICE_HAS_ext_oneapi_bindless_sampled_image_fetch_3d__ 0
 #endif
+
+#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_queue_profiling_tag__
+// __SYCL_ASPECT(ext_oneapi_queue_profiling_tag, 73)
+#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_queue_profiling_tag__ 0
+#endif

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

@@ -0,0 +1,44 @@
+//==--------- profiling_tag.hpp --- SYCL profiling tag extension -----------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <sycl/aspects.hpp>
+#include <sycl/event.hpp>
+#include <sycl/handler.hpp>
+#include <sycl/properties/queue_properties.hpp>
+#include <sycl/queue.hpp>
+
+namespace sycl {
+inline namespace _V1 {
+namespace ext::oneapi::experimental {
+
+inline event submit_profiling_tag(queue &Queue,
+                                  const sycl::detail::code_location &CodeLoc =
+                                      sycl::detail::code_location::current()) {
+  if (Queue.get_device().has(aspect::ext_oneapi_queue_profiling_tag)) {
+    return Queue.submit(
+        [=](handler &CGH) {
+          sycl::detail::HandlerAccess::internalProfilingTagImpl(CGH);
+        },
+        CodeLoc);
+  }
+
+  // If it is not supported natively on the device, we use another path if
+  // profiling is enabled.
+  if (!Queue.has_property<sycl::property::queue::enable_profiling>())
+    throw sycl::exception(
+        make_error_code(errc::invalid),
+        "Device must either have aspect::ext_oneapi_queue_profiling_tag or the "
+        "queue must have profiling enabled.");
+  return Queue.ext_oneapi_submit_barrier();
+}
+
+} // namespace ext::oneapi::experimental
+} // namespace _V1
+} // namespace sycl

sycl/include/sycl/handler.hpp

@@ -175,6 +175,7 @@ class stream_impl;
 template <typename DataT, int Dimensions, access::mode AccessMode,
           access::target AccessTarget, access::placeholder IsPlaceholder>
 class image_accessor;
+class HandlerAccess;
 template <typename RetType, typename Func, typename Arg>
 static Arg member_ptr_helper(RetType (Func::*)(Arg) const);
 
@@ -3676,12 +3677,29 @@ class __SYCL_EXPORT handler {
   // Set that an ND Range was used during a call to parallel_for
   void setNDRangeUsed(bool Value);
 
+  inline void internalProfilingTagImpl() {
+    throwIfActionIsCreated();
+    setType(detail::CG::ProfilingTag);
+  }
+
+  friend class detail::HandlerAccess;
+
 protected:
   /// Registers event dependencies in this command group.
   void depends_on(const detail::EventImplPtr &Event);
   /// Registers event dependencies in this command group.
   void depends_on(const std::vector<detail::EventImplPtr> &Events);
 };
+
+namespace detail {
+class HandlerAccess {
+public:
+  static void internalProfilingTagImpl(handler &Handler) {
+    Handler.internalProfilingTagImpl();
+  }
+};
+} // namespace detail
+
 } // namespace _V1
 } // namespace sycl
 

sycl/include/sycl/info/aspects.def

@@ -67,3 +67,4 @@ __SYCL_ASPECT(ext_oneapi_bindless_sampled_image_fetch_2d_usm, 69)
 __SYCL_ASPECT(ext_oneapi_bindless_sampled_image_fetch_2d, 70)
 __SYCL_ASPECT(ext_oneapi_bindless_sampled_image_fetch_3d_usm, 71)
 __SYCL_ASPECT(ext_oneapi_bindless_sampled_image_fetch_3d, 72)
+__SYCL_ASPECT(ext_oneapi_queue_profiling_tag, 73)

sycl/include/sycl/sycl.hpp

@@ -96,6 +96,7 @@
 #include <sycl/ext/oneapi/experimental/group_load_store.hpp>
 #include <sycl/ext/oneapi/experimental/opportunistic_group.hpp>
 #include <sycl/ext/oneapi/experimental/prefetch.hpp>
+#include <sycl/ext/oneapi/experimental/profiling_tag.hpp>
 #include <sycl/ext/oneapi/experimental/root_group.hpp>
 #include <sycl/ext/oneapi/experimental/tangle_group.hpp>
 #include <sycl/ext/oneapi/filter_selector.hpp>

sycl/plugins/cuda/pi_cuda.cpp

@@ -631,6 +631,14 @@ pi_result piextEventCreateWithNativeHandle(pi_native_handle NativeHandle,
                                                  OwnNativeHandle, Event);
 }
 
+pi_result piEnqueueTimestampRecordingExp(pi_queue Queue, pi_bool Blocking,
+                                         pi_uint32 NumEventsInWaitList,
+                                         const pi_event *EventWaitList,
+                                         pi_event *Event) {
+  return pi2ur::piEnqueueTimestampRecordingExp(
+      Queue, Blocking, NumEventsInWaitList, EventWaitList, Event);
+}
+
 pi_result piSamplerCreate(pi_context Context,
                           const pi_sampler_properties *SamplerProperties,
                           pi_sampler *RetSampler) {