Added the profiling stream for CUDA backend

Author: keyradical

source/adapters/cuda/event.cpp

@@ -55,8 +55,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(cuEventRecord(EvQueued, 0));
+      UR_CHECK_ERROR(cuEventRecord(EvQueued, Queue->getProfilingStream()));
       UR_CHECK_ERROR(cuEventRecord(EvStart, Stream));
     }
   } catch (ur_result_t Err) {

source/adapters/cuda/event.hpp

@@ -90,6 +90,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(cuEventCreate(
         &EvEnd, RequiresTimings ? CU_EVENT_DEFAULT : CU_EVENT_DISABLE_TIMING));

source/adapters/cuda/queue.cpp

@@ -201,6 +201,11 @@ UR_APIEXPORT ur_result_t UR_APICALL urQueueRelease(ur_queue_handle_t hQueue) {
       UR_CHECK_ERROR(cuStreamDestroy(S));
     });
 
+    if (hQueue->IsProfStreamCreated) {
+      UR_CHECK_ERROR(cuStreamSynchronize(hQueue->getProfilingStream()));
+      UR_CHECK_ERROR(cuStreamDestroy(hQueue->getProfilingStream()));
+    }
+
     return UR_RESULT_SUCCESS;
   } catch (ur_result_t Err) {
     return Err;

source/adapters/cuda/queue.hpp

@@ -9,10 +9,12 @@
 //===----------------------------------------------------------------------===//
 #pragma once
 
+#include "common.hpp"
 #include <ur/ur.hpp>
 
 #include <algorithm>
 #include <cuda.h>
+#include <mutex>
 #include <vector>
 
 using ur_stream_guard_ = std::unique_lock<std::mutex>;
@@ -27,6 +29,9 @@ 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;
+  bool IsProfStreamCreated{false};
   // delay_compute_ 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
@@ -64,8 +69,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},
@@ -99,6 +104,19 @@ struct ur_queue_handle_t_ {
   native_type get() { return getNextComputeStream(); };
   ur_device_handle_t getDevice() const noexcept { return Device; };
 
+  // Function which creates the profiling stream. Called only if profiling is
+  // enabled.
+  void createProfilingStream() {
+    static std::once_flag ProfStreamFlag;
+    std::call_once(ProfStreamFlag, [&]() {
+      UR_CHECK_ERROR(
+          cuStreamCreateWithPriority(&ProfStream, CU_STREAM_NON_BLOCKING, 0));
+      IsProfStreamCreated = 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()) {

source/adapters/hip/event.cpp

@@ -21,8 +21,8 @@ ur_event_handle_t_::ur_event_handle_t_(ur_command_t Type,
                                        uint32_t StreamToken)
     : CommandType{Type}, RefCount{1}, HasOwnership{true},
       HasBeenWaitedOn{false}, IsRecorded{false}, IsStarted{false},
-      StreamToken{StreamToken}, EventId{0}, EvEnd{EvEnd}, EvQueued{EvQueued},
-      EvStart{EvStart}, Queue{Queue}, Stream{Stream}, Context{Context} {
+      StreamToken{StreamToken}, EventId{0}, EvEnd{EvEnd}, EvStart{EvStart},
+      EvQueued{EvQueued}, Queue{Queue}, Stream{Stream}, Context{Context} {
   urQueueRetain(Queue);
   urContextRetain(Context);
 }
@@ -50,7 +50,6 @@ 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, Queue->getProfilingStream()));
       UR_CHECK_ERROR(hipEventRecord(EvStart, Stream));
     }

source/adapters/hip/event.hpp

@@ -112,8 +112,9 @@ struct ur_event_handle_t_ {
   // This constructor is private to force programmers to use the makeNative /
   // make_user static members in order to create a ur_event_handle_t for HIP.
   ur_event_handle_t_(ur_command_t Type, ur_context_handle_t Context,
-                     ur_queue_handle_t Queue, hipStream_t Stream,
-                     uint32_t StreamToken);
+                     ur_queue_handle_t Queue, native_type EvEnd,
+                     native_type EvQueued, native_type EvStart,
+                     hipStream_t Stream, uint32_t StreamToken);
 
   // This constructor is private to force programmers to use the
   // makeWithNative for event interop

source/adapters/hip/queue.cpp

@@ -222,9 +222,9 @@ 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));
+    if (hQueue->IsProfStreamCreated) {
+      UR_CHECK_ERROR(hipStreamSynchronize(hQueue->getProfilingStream()));
+      UR_CHECK_ERROR(hipStreamDestroy(hQueue->getProfilingStream()));
     }
 
     return UR_RESULT_SUCCESS;

source/adapters/hip/queue.hpp

@@ -10,6 +10,9 @@
 #pragma once
 
 #include "common.hpp"
+#include <hip/hip_runtime.h>
+#include <mutex>
+#include <vector>
 
 using ur_stream_quard = std::unique_lock<std::mutex>;
 
@@ -24,8 +27,7 @@ struct ur_queue_handle_t_ {
   std::vector<native_type> TransferStreams;
   // Stream used solely when profiling is enabled
   native_type ProfStream;
-  static std::once_flag ProfStreamFlag;
-  bool ProfStreamCreated;
+  bool IsProfStreamCreated{false};
   // 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
@@ -99,13 +101,14 @@ 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.
+  // Function which creates the profiling stream. Called only if profiling is
+  // enabled.
   void createProfilingStream() {
-    std::call_once(ProfStreamFlag, []() {
+    static std::once_flag ProfStreamFlag;
+    std::call_once(ProfStreamFlag, [&]() {
       UR_CHECK_ERROR(
-          hipEventCreateWithFlags(&ProfStream, hipStreamNonBlocking));
-      ProfStreamCreated = true;
+          hipStreamCreateWithFlags(&ProfStream, hipStreamNonBlocking));
+      IsProfStreamCreated = true;
     });
   }
   native_type getProfilingStream() { return ProfStream; }