SWDEV-429053 - Add check for StreamLegacy

Signed-off-by: sdashmiz <shadi.dashmiz@amd.com>
Change-Id: I402185a3b81935aaa1c8c4963407b8de21c72d8a

hipamd/src/hip_context.cpp

@@ -91,7 +91,7 @@ void setCurrentDevice(unsigned int index) {
 }
 
 hip::Stream* getStream(hipStream_t stream, bool wait) {
- if (stream == nullptr) {
+ if (stream == nullptr || stream == hipStreamLegacy) {
     return getNullStream(wait);
   } else {
     hip::Stream* hip_stream = reinterpret_cast<hip::Stream*>(stream);

hipamd/src/hip_graph.cpp

@@ -855,10 +855,11 @@ hipError_t hipStreamIsCapturing_common(hipStream_t stream, hipStreamCaptureStatu
   if (!hip::isValid(stream)) {
     return hipErrorContextIsDestroyed;
   }
-  if (hip::Stream::StreamCaptureBlocking() == true && stream == nullptr) {
+  if (hip::Stream::StreamCaptureBlocking() == true &&
+      (stream == nullptr || stream == hipStreamLegacy)) {
     return hipErrorStreamCaptureImplicit;
   }
-  if (stream == nullptr) {
+  if (stream == nullptr || stream == hipStreamLegacy) {
     *pCaptureStatus = hipStreamCaptureStatusNone;
   } else {
     *pCaptureStatus = reinterpret_cast<hip::Stream*>(stream)->GetCaptureStatus();
@@ -898,7 +899,7 @@ hipError_t hipStreamBeginCapture_common(hipStream_t stream, hipStreamCaptureMode
     return hipErrorContextIsDestroyed;
   }
   // capture cannot be initiated on legacy stream
-  if (stream == nullptr) {
+  if (stream == nullptr || stream == hipStreamLegacy) {
     return hipErrorStreamCaptureUnsupported;
   }
   if (mode < hipStreamCaptureModeGlobal || mode > hipStreamCaptureModeRelaxed) {
@@ -977,7 +978,7 @@ hipError_t hipStreamEndCapture_common(hipStream_t stream, hip::Graph** pGraph) {
   if (pGraph == nullptr) {
     return hipErrorInvalidValue;
   }
-  if (stream == nullptr) {
+  if (stream == nullptr || stream == hipStreamLegacy) {
     return hipErrorIllegalState;
   }
   if (!hip::isValid(stream)) {
@@ -1771,10 +1772,11 @@ hipError_t hipStreamGetCaptureInfo_common(hipStream_t stream,
   if (!hip::isValid(stream)) {
     return hipErrorContextIsDestroyed;
   }
-  if (hip::Stream::StreamCaptureBlocking() == true && stream == nullptr) {
+  if (hip::Stream::StreamCaptureBlocking() == true &&
+      (stream == nullptr || stream == hipStreamLegacy)) {
     return hipErrorStreamCaptureImplicit;
   }
-  if (stream == nullptr) {
+  if (stream == nullptr || stream == hipStreamLegacy) {
     *pCaptureStatus = hipStreamCaptureStatusNone;
     return hipSuccess;
   }
@@ -1807,10 +1809,11 @@ hipError_t hipStreamGetCaptureInfo_v2_common(hipStream_t stream,
   if (captureStatus_out == nullptr) {
     return hipErrorInvalidValue;
   }
-  if (hip::Stream::StreamCaptureBlocking() == true && stream == nullptr) {
+  if (hip::Stream::StreamCaptureBlocking() == true &&
+      (stream == nullptr || stream == hipStreamLegacy)) {
     return hipErrorStreamCaptureImplicit;
   }
-  if (stream == nullptr) {
+  if (stream == nullptr || stream == hipStreamLegacy) {
     *captureStatus_out = hipStreamCaptureStatusNone;
     return hipSuccess;
   }

hipamd/src/hip_hmm.cpp

@@ -108,10 +108,12 @@ hipError_t hipMemPrefetchAsync(const void* dev_ptr, size_t count, int device,
   // Pick the specified stream or Null one from the provided device
   if (device == hipCpuDeviceId) {
     cpu_access = true;
-    hip_stream = (stream == nullptr) ? hip::getCurrentDevice()->NullStream() : hip::getStream(stream);
+    hip_stream = (stream == nullptr || stream == hipStreamLegacy) ?
+                  hip::getCurrentDevice()->NullStream() : hip::getStream(stream);
   } else {
     dev = g_devices[device]->devices()[0];
-    hip_stream = (stream == nullptr) ? g_devices[device]->NullStream() : hip::getStream(stream);
+    hip_stream = (stream == nullptr || stream == hipStreamLegacy) ?
+                  g_devices[device]->NullStream() : hip::getStream(stream);
   }
 
   if (hip_stream == nullptr) {
@@ -250,8 +252,8 @@ hipError_t hipStreamAttachMemAsync(hipStream_t stream, void* dev_ptr,
   // host-accessible region of system-allocated pageable memory.
   // This type of memory may only be specified if the device associated with the
   // stream reports a non-zero value for the device attribute hipDevAttrPageableMemoryAccess.
-  hip::Stream* hip_stream = (stream == nullptr) ? hip::getCurrentDevice()->NullStream()
-                                                : hip::getStream(stream);
+  hip::Stream* hip_stream = (stream == nullptr || stream == hipStreamLegacy) ?
+                             hip::getCurrentDevice()->NullStream() : hip::getStream(stream);
   size_t offset = 0;
   amd::Memory* memObj = getMemoryObject(dev_ptr, offset);
   if (memObj == nullptr) {

hipamd/src/hip_internal.hpp

@@ -230,7 +230,7 @@ const char* ihipGetErrorName(hipError_t hip_error);
 
 #define STREAM_CAPTURE(name, stream, ...)                                                          \
   hip::getStreamPerThread(stream);                                                                 \
-  if (stream != nullptr &&                                                                         \
+  if (stream != nullptr && stream != hipStreamLegacy &&                                            \
       reinterpret_cast<hip::Stream*>(stream)->GetCaptureStatus() ==                                \
           hipStreamCaptureStatusActive) {                                                          \
     hipError_t status = hip::capture##name(stream, ##__VA_ARGS__);                                 \
@@ -242,7 +242,7 @@ const char* ihipGetErrorName(hipError_t hip_error);
   }
 
 #define PER_THREAD_DEFAULT_STREAM(stream)                                                         \
-  if (stream == nullptr) {                                                                        \
+  if (stream == nullptr || stream == hipStreamLegacy) {                                           \
     stream = getPerThreadDefaultStream();                                                         \
   }
 

hipamd/src/hip_memory.cpp

@@ -674,7 +674,7 @@ hipError_t hipMemcpy_common(void* dst, const void* src, size_t sizeBytes,
   CHECK_STREAM_CAPTURING();
   hip::Stream* hip_stream = nullptr;
 
-  if (stream != nullptr) {
+  if (stream != nullptr && stream != hipStreamLegacy) {
     hip_stream = hip::getStream(stream);
   } else {
     hip_stream = hip::getNullStream();

hipamd/src/hip_mempool.cpp

@@ -92,8 +92,8 @@ hipError_t hipMallocAsync(void** dev_ptr, size_t size, hipStream_t stream) {
     *dev_ptr = nullptr;
     HIP_RETURN(hipSuccess);
   }
-  auto hip_stream = (stream == nullptr) ? hip::getCurrentDevice()->NullStream() :
-    reinterpret_cast<hip::Stream*>(stream);
+  auto hip_stream = (stream == nullptr || stream == hipStreamLegacy) ?
+    hip::getCurrentDevice()->NullStream() : reinterpret_cast<hip::Stream*>(stream);
   auto device = hip_stream->GetDevice();
   auto mem_pool = device->GetCurrentMemoryPool();
 
@@ -147,8 +147,8 @@ hipError_t hipFreeAsync(void* dev_ptr, hipStream_t stream) {
 
   STREAM_CAPTURE(hipFreeAsync, stream, dev_ptr);
 
-  auto hip_stream = (stream == nullptr) ? hip::getCurrentDevice()->NullStream()
-                    : reinterpret_cast<hip::Stream*>(stream);
+  auto hip_stream = (stream == nullptr || stream == hipStreamLegacy) ?
+    hip::getCurrentDevice()->NullStream(): reinterpret_cast<hip::Stream*>(stream);
 
   hip::Event* event = nullptr;
   bool graph_in_use = false;
@@ -192,7 +192,7 @@ hipError_t hipFreeAsync(void* dev_ptr, hipStream_t stream) {
         }
       }
     }
-    
+
     auto cmd = new FreeAsyncCommand(*hip_stream, dev_ptr, event);
     if (cmd == nullptr) {
       HIP_RETURN(hipErrorUnknown);
@@ -367,8 +367,8 @@ hipError_t hipMallocFromPoolAsync(
   STREAM_CAPTURE(hipMallocAsync, stream, mem_pool, size, dev_ptr);
 
   auto mpool = reinterpret_cast<hip::MemoryPool*>(mem_pool);
-  auto hip_stream = (stream == nullptr) ? hip::getCurrentDevice()->NullStream() :
-    reinterpret_cast<hip::Stream*>(stream);
+  auto hip_stream = (stream == nullptr || stream == hipStreamLegacy) ?
+    hip::getCurrentDevice()->NullStream() : reinterpret_cast<hip::Stream*>(stream);
   *dev_ptr = mpool->AllocateMemory(size, hip_stream);
   if (*dev_ptr == nullptr) {
     HIP_RETURN(hipErrorOutOfMemory);

hipamd/src/hip_stream.cpp

@@ -80,7 +80,7 @@ void Stream::Destroy(hip::Stream* stream) {
 // ================================================================================================
 bool isValid(hipStream_t& stream) {
   // NULL stream is always valid
-  if (stream == nullptr) {
+  if (stream == nullptr || stream == hipStreamLegacy) {
     return true;
   }