Added a profiling stream to the Queue

Author: Konrad Kusiak

source/adapters/hip/event.cpp

@@ -51,7 +51,7 @@ ur_result_t ur_event_handle_t_::start() {
   try {
     if (Queue->URFlags & UR_QUEUE_FLAG_PROFILING_ENABLE || isTimestampEvent()) {
       // NOTE: This relies on the default stream to be unused.
-      UR_CHECK_ERROR(hipEventRecord(EvQueued, 0));
+      UR_CHECK_ERROR(hipEventRecord(EvQueued, Queue->getProfilingStream()));
       UR_CHECK_ERROR(hipEventRecord(EvStart, Stream));
     }
   } catch (ur_result_t Error) {

source/adapters/hip/event.hpp

@@ -83,6 +83,9 @@ struct ur_event_handle_t_ {
     const bool RequiresTimings =
         Queue->URFlags & UR_QUEUE_FLAG_PROFILING_ENABLE ||
         Type == UR_COMMAND_TIMESTAMP_RECORDING_EXP;
+    if (RequiresTimings) {
+      Queue->createProfilingStream();
+    }
     native_type EvEnd{nullptr}, EvQueued{nullptr}, EvStart{nullptr};
     UR_CHECK_ERROR(hipEventCreateWithFlags(
         &EvEnd, RequiresTimings ? hipEventDefault : hipEventDisableTiming));

source/adapters/hip/queue.cpp

@@ -222,6 +222,11 @@ UR_APIEXPORT ur_result_t UR_APICALL urQueueRelease(ur_queue_handle_t hQueue) {
       UR_CHECK_ERROR(hipStreamDestroy(S));
     });
 
+    if (hQueue->ProfStreamCreated) {
+      UR_CHECK_ERROR(hipStreamSynchronize(ProfStream));
+      UR_CHECK_ERROR(hipStreamDestroy(ProfStream));
+    }
+
     return UR_RESULT_SUCCESS;
   } catch (ur_result_t Err) {
     return Err;

source/adapters/hip/queue.hpp

@@ -22,6 +22,10 @@ struct ur_queue_handle_t_ {
 
   std::vector<native_type> ComputeStreams;
   std::vector<native_type> TransferStreams;
+  // Stream used solely when profiling is enabled
+  native_type ProfStream;
+  static std::once_flag ProfStreamFlag;
+  bool ProfStreamCreated;
   // DelayCompute keeps track of which streams have been recently reused and
   // their next use should be delayed. If a stream has been recently reused it
   // will be skipped the next time it would be selected round-robin style. When
@@ -59,8 +63,8 @@ struct ur_queue_handle_t_ {
                      ur_context_handle_t Context, ur_device_handle_t Device,
                      unsigned int Flags, ur_queue_flags_t URFlags, int Priority,
                      bool BackendOwns = true)
-      : ComputeStreams{std::move(ComputeStreams)}, TransferStreams{std::move(
-                                                       TransferStreams)},
+      : ComputeStreams{std::move(ComputeStreams)},
+        TransferStreams{std::move(TransferStreams)},
         DelayCompute(this->ComputeStreams.size(), false),
         ComputeAppliedBarrier(this->ComputeStreams.size()),
         TransferAppliedBarrier(this->TransferStreams.size()), Context{Context},
@@ -95,6 +99,17 @@ struct ur_queue_handle_t_ {
   native_type getNextTransferStream();
   native_type get() { return getNextComputeStream(); };
 
+  // Function which creates the profiling stream. Called only from makeNative
+  // in event handle, if the profiling is enabled.
+  void createProfilingStream() {
+    std::call_once(ProfStreamFlag, []() {
+      UR_CHECK_ERROR(
+          hipEventCreateWithFlags(&ProfStream, hipStreamNonBlocking));
+      ProfStreamCreated = true;
+    });
+  }
+  native_type getProfilingStream() { return ProfStream; }
+
   bool hasBeenSynchronized(uint32_t StreamToken) {
     // stream token not associated with one of the compute streams
     if (StreamToken == std::numeric_limits<uint32_t>::max()) {