SWDEV-292547 - hipStreamPerThread support

Change-Id: Id621ce073b0fee9eac03c59ffb78b197fda4ddb5


[ROCm/clr commit: c3ca1faee7]

projects/clr/hipamd/src/hip_stream.cpp

@@ -172,12 +172,16 @@ void Stream::destroyAllStreams(int deviceId) {
 }
 
 // ================================================================================================
-bool isValid(hipStream_t stream) {
+bool isValid(hipStream_t& stream) {
   // NULL stream is always valid
   if (stream == nullptr) {
     return true;
   }
 
+  if (hipStreamPerThread == stream) {
+    getStreamPerThread(stream);
+  }
+
   hip::Stream* s = reinterpret_cast<hip::Stream*>(stream);
   amd::ScopedLock lock(streamSetLock);
   if (streamSet.find(s) == streamSet.end()) {
@@ -266,6 +270,36 @@ static hipError_t ihipStreamCreate(hipStream_t* stream,
   return hipSuccess;
 }
 
+// ================================================================================================
+class stream_per_thread {
+private:
+  hipStream_t m_stream;
+public:
+  stream_per_thread():m_stream(nullptr) {}
+  stream_per_thread(const stream_per_thread& ) = delete;
+  void operator=(const stream_per_thread& ) = delete;
+  ~stream_per_thread() {
+    if (m_stream != nullptr && hip::isValid(m_stream)) {
+      delete reinterpret_cast<hip::Stream*>(m_stream);
+    }
+  }
+
+  hipStream_t& get() {
+    if (m_stream == nullptr) {
+      ihipStreamCreate(&m_stream, hipStreamDefault, hip::Stream::Priority::Normal);
+    }
+    return m_stream;
+  }
+};
+thread_local stream_per_thread streamPerThreadObj;
+
+// ================================================================================================
+hipError_t getStreamPerThread(hipStream_t& stream) {
+  if (stream == hipStreamPerThread) {
+    stream = streamPerThreadObj.get();
+  }
+}
+
 // ================================================================================================
 hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags) {
   HIP_INIT_API(hipStreamCreateWithFlags, stream, flags);
@@ -520,6 +554,7 @@ hipError_t hipExtStreamCreateWithCUMask(hipStream_t* stream, uint32_t cuMaskSize
 // ================================================================================================
 hipError_t hipStreamGetPriority(hipStream_t stream, int* priority) {
   HIP_INIT_API(hipStreamGetPriority, stream, priority);
+
   if ((priority != nullptr) && (stream != nullptr)) {
     if (!hip::isValid(stream)) {
       return HIP_RETURN(hipErrorContextIsDestroyed);
@@ -576,7 +611,7 @@ hipError_t hipExtStreamGetCUMask(hipStream_t stream, uint32_t cuMaskSize, uint32
 
   // if the stream is null then either return globalCUMask_ (if it is defined)
   // or return defaultCUMask
-  if (stream == nullptr) {
+  if (stream == nullptr || stream == hipStreamPerThread) {
     if (info.globalCUMask_.size() != 0) {
       std::copy(info.globalCUMask_.begin(), info.globalCUMask_.end(), cuMask);
     } else {