SWDEV-527781 - Remove Stream Validation in HIP APIs

[ROCm/clr commit: 0457b634f8]

projects/clr/hipamd/src/hip_event.cpp

@@ -390,9 +390,6 @@ hipError_t hipEventRecord_common(hipEvent_t event, hipStream_t stream, unsigned
     return hipErrorInvalidHandle;
   }
   getStreamPerThread(stream);
-  if (!hip::isValid(stream)) {
-    return hipErrorContextIsDestroyed;
-  }
   hip::Event* e = reinterpret_cast<hip::Event*>(event);
   hip::Stream* s = reinterpret_cast<hip::Stream*>(stream);
   hip::Stream* hip_stream = hip::getStream(stream);

projects/clr/hipamd/src/hip_graph.cpp

@@ -42,9 +42,7 @@ inline hipError_t ihipGraphUpload(hipGraphExec_t graphExec, hipStream_t stream)
   if (graphExec == nullptr) {
     return hipErrorInvalidValue;
   }
-  if (!hip::isValid(stream)) {
-    return hipErrorContextIsDestroyed;
-  }
+  getStreamPerThread(stream);
   if (!hip::GraphExec::isGraphExecValid(reinterpret_cast<hip::GraphExec*>(graphExec))) {
     return hipErrorInvalidValue;
   }
@@ -1026,9 +1024,7 @@ hipError_t hipStreamIsCapturing_common(hipStream_t stream, hipStreamCaptureStatu
   if (pCaptureStatus == nullptr) {
     return hipErrorInvalidValue;
   }
-  if (!hip::isValid(stream)) {
-    return hipErrorContextIsDestroyed;
-  }
+  getStreamPerThread(stream);
   if (hip::Stream::StreamCaptureBlocking() == true &&
       (stream == nullptr || stream == hipStreamLegacy)) {
     return hipErrorStreamCaptureImplicit;
@@ -1069,9 +1065,7 @@ hipError_t hipThreadExchangeStreamCaptureMode(hipStreamCaptureMode* mode) {
 
 hipError_t hipStreamBeginCapture_common(hipStream_t stream, hipStreamCaptureMode mode,
                                         hipGraph_t graph = nullptr) {
-  if (!hip::isValid(stream)) {
-    return hipErrorContextIsDestroyed;
-  }
+  getStreamPerThread(stream);
   // capture cannot be initiated on legacy stream
   if (stream == nullptr || stream == hipStreamLegacy) {
     return hipErrorStreamCaptureUnsupported;
@@ -1591,9 +1585,7 @@ hipError_t hipGraphExecDestroy(hipGraphExec_t pGraphExec) {
 }
 
 hipError_t ihipGraphLaunch(hip::GraphExec* graphExec, hipStream_t stream) {
-  if (!hip::isValid(stream)) {
-    return hipErrorContextIsDestroyed;
-  }
+  getStreamPerThread(stream);
   hip::Stream* launch_stream = hip::getStream(stream);
   return graphExec->Run(launch_stream);
 }
@@ -1605,9 +1597,6 @@ hipError_t hipGraphLaunch_common(hip::GraphExec* graphExec, hipStream_t stream)
   if (graphExec->GetNodeCount() == 0) {
     return hipSuccess;
   }
-  if (!hip::isValid(stream)) {
-    return hipErrorContextIsDestroyed;
-  }
   return ihipGraphLaunch(graphExec, stream);
 }
 
@@ -2004,9 +1993,7 @@ hipError_t hipStreamGetCaptureInfo_common(hipStream_t stream,
   if (pCaptureStatus == nullptr) {
     return hipErrorInvalidValue;
   }
-  if (!hip::isValid(stream)) {
-    return hipErrorContextIsDestroyed;
-  }
+  getStreamPerThread(stream);
   if (hip::Stream::StreamCaptureBlocking() == true &&
       (stream == nullptr || stream == hipStreamLegacy)) {
     return hipErrorStreamCaptureImplicit;

projects/clr/hipamd/src/hip_hmm.cpp

@@ -82,9 +82,7 @@ hipError_t hipMemPrefetchAsync(const void* dev_ptr, size_t count, int device,
     HIP_RETURN(hipErrorInvalidValue);
   }
 
-  if (!hip::isValid(stream)) {
-    HIP_RETURN(hipErrorContextIsDestroyed);
-  }
+  getStreamPerThread(stream);
 
   size_t offset = 0;
   amd::Memory* memObj = getMemoryObject(dev_ptr, offset);
@@ -238,9 +236,7 @@ hipError_t hipStreamAttachMemAsync(hipStream_t stream, void* dev_ptr,
     HIP_RETURN(hipErrorInvalidValue);
   }
 
-  if (!hip::isValid(stream)) {
-    HIP_RETURN(hipErrorContextIsDestroyed);
-  }
+  getStreamPerThread(stream);
 
   if (flags != hipMemAttachGlobal && flags != hipMemAttachHost && flags != hipMemAttachSingle) {
     HIP_RETURN(hipErrorInvalidValue);

projects/clr/hipamd/src/hip_memory.cpp

@@ -1430,13 +1430,11 @@ hipError_t hipMemcpyAsync_common(void* dst, const void* src, size_t sizeBytes,
   if (static_cast<uint32_t>(kind) > hipMemcpyDefault && kind != hipMemcpyDeviceToDeviceNoCU) {
     return hipErrorInvalidMemcpyDirection;
   }
+  getStreamPerThread(stream);
   hip::Stream* hip_stream = hip::getStream(stream);
   if (hip_stream == nullptr) {
     return hipErrorInvalidValue;
   }
-  if (!hip::isValid(stream)) {
-    return hipErrorContextIsDestroyed;
-  }
   return ihipMemcpy(dst, src, sizeBytes, kind, *hip_stream, true);
 }
 
@@ -2360,9 +2358,7 @@ hipError_t ihipMemcpyParam3D(const HIP_MEMCPY3D* pCopy, hipStream_t stream, bool
   if (pCopy == nullptr) {
     return hipErrorInvalidValue;
   }
-  if (!hip::isValid(stream)) {
-    return hipErrorContextIsDestroyed;
-  }
+  getStreamPerThread(stream);
   hipMemoryType srcMemoryType;
   hipMemoryType dstMemoryType;
   ihipCopyMemParamSet(pCopy, srcMemoryType, dstMemoryType);
@@ -2448,9 +2444,7 @@ hipError_t hipMemcpy2DValidateParams(hipMemcpyKind kind, hipStream_t stream = nu
     return hipErrorInvalidMemcpyDirection;
   }
 
-  if (!hip::isValid(stream)) {
-    return hipErrorInvalidValue;
-  }
+  getStreamPerThread(stream);
 
   return hipSuccess;
 }

projects/clr/hipamd/src/hip_mempool.cpp

@@ -85,9 +85,7 @@ hipError_t hipMallocAsync(void** dev_ptr, size_t size, hipStream_t stream) {
   if (dev_ptr == nullptr) {
     HIP_RETURN(hipErrorInvalidValue);
   }
-  if (!hip::isValid(stream)) {
-    HIP_RETURN(hipErrorInvalidHandle);
-  }
+  getStreamPerThread(stream);
   if (size == 0) {
     *dev_ptr = nullptr;
     HIP_RETURN(hipSuccess);
@@ -147,9 +145,7 @@ class FreeAsyncCommand : public amd::Command {
 hipError_t hipFreeAsync(void* dev_ptr, hipStream_t stream) {
   HIP_INIT_API(hipFreeAsync, dev_ptr, stream);
 
-  if (!hip::isValid(stream)) {
-    HIP_RETURN(hipErrorInvalidHandle);
-  }
+  getStreamPerThread(stream);
 
   hip::Stream* s = reinterpret_cast<hip::Stream*>(stream);
   auto hip_stream = (stream == nullptr || stream == hipStreamLegacy) ?
@@ -376,9 +372,7 @@ hipError_t hipMallocFromPoolAsync(
   if ((dev_ptr == nullptr) || (mem_pool == nullptr)) {
     HIP_RETURN(hipErrorInvalidValue);
   }
-  if (!hip::isValid(stream)) {
-    HIP_RETURN(hipErrorInvalidHandle);
-  }
+  getStreamPerThread(stream);
   if (size == 0) {
     *dev_ptr = nullptr;
     HIP_RETURN(hipSuccess);

projects/clr/hipamd/src/hip_peer.cpp

@@ -209,9 +209,7 @@ hipError_t hipMemcpyPeerAsync(void* dst, int dstDevice, const void* src, int src
       srcDevice < 0 || dstDevice < 0) {
     HIP_RETURN(hipErrorInvalidDevice);
   }
-  if (!hip::isValid(stream)) {
-    return hipErrorContextIsDestroyed;
-  }
+  getStreamPerThread(stream);
   hip::Stream* hip_stream = hip::getStream(stream);
   if (hip_stream == nullptr) {
     return hipErrorInvalidValue;

projects/clr/hipamd/src/hip_stream.cpp

@@ -323,9 +323,7 @@ hipError_t hipDeviceGetStreamPriorityRange(int* leastPriority, int* greatestPrio
 // ================================================================================================
 hipError_t hipStreamGetFlags_common(hipStream_t stream, unsigned int* flags) {
   if ((flags != nullptr) && (stream != nullptr)) {
-    if (!hip::isValid(stream)) {
-      return hipErrorContextIsDestroyed;
-    }
+    getStreamPerThread(stream);
     *flags = reinterpret_cast<hip::Stream*>(stream)->Flags();
   } else {
     return hipErrorInvalidValue;
@@ -349,9 +347,7 @@ hipError_t hipStreamGetFlags_spt(hipStream_t stream, unsigned int* flags) {
 
 // ================================================================================================
 hipError_t hipStreamSynchronize_common(hipStream_t stream) {
-  if (!hip::isValid(stream)) {
-    HIP_RETURN(hipErrorContextIsDestroyed);
-  }
+  getStreamPerThread(stream);
   if (stream != nullptr && stream != hipStreamLegacy) {
     // If still capturing return error
     if (hip::Stream::StreamCaptureOngoing(stream) == true) {
@@ -398,9 +394,6 @@ hipError_t hipStreamDestroy(hipStream_t stream) {
   if (stream == hipStreamPerThread || stream == hipStreamLegacy) {
     HIP_RETURN(hipErrorInvalidResourceHandle);
   }
-  if (!hip::isValid(stream)) {
-    HIP_RETURN(hipErrorContextIsDestroyed);
-  }
   hip::Stream* s = reinterpret_cast<hip::Stream*>(stream);
   if (s->GetCaptureStatus() != hipStreamCaptureStatusNone) {
     if (s->GetParentStream() != nullptr) {
@@ -448,9 +441,10 @@ void WaitThenDecrementSignal(hipStream_t stream, hipError_t status, void* user_d
 // ================================================================================================
 hipError_t hipStreamWaitEvent_common(hipStream_t stream, hipEvent_t event, unsigned int flags) {
   hipError_t status = hipSuccess;
-  if (event == nullptr || !hip::isValid(stream)) {
+  if (event == nullptr) {
     return hipErrorInvalidHandle;
   }
+  getStreamPerThread(stream);
   hip::Stream* waitStream = hip::getStream(stream);
   hip::Event* e = reinterpret_cast<hip::Event*>(event);
   auto eventStreamHandle = reinterpret_cast<hipStream_t>(e->GetCaptureStream());
@@ -511,9 +505,7 @@ hipError_t hipStreamWaitEvent_spt(hipStream_t stream, hipEvent_t event, unsigned
 
 // ================================================================================================
 hipError_t hipStreamQuery_common(hipStream_t stream) {
-  if (!hip::isValid(stream)) {
-    return hipErrorContextIsDestroyed;
-  }
+  getStreamPerThread(stream);
   if (stream != nullptr) {
     // If still capturing return error
     if (hip::Stream::StreamCaptureOngoing(stream) == true) {
@@ -566,10 +558,7 @@ hipError_t hipStreamQuery_spt(hipStream_t stream) {
 }
 
 hipError_t streamCallback_common(hipStream_t stream, StreamCallback* cbo, void* userData) {
-  if (!hip::isValid(stream)) {
-    return hipErrorContextIsDestroyed;
-  }
-
+  getStreamPerThread(stream);
   hip::Stream* hip_stream = hip::getStream(stream);
   amd::Command* last_command = hip_stream->getLastQueuedCommand(true);
   amd::Command::EventWaitList eventWaitList;
@@ -688,9 +677,7 @@ hipError_t hipStreamGetPriority_common(hipStream_t stream, int* priority) {
   }
 
   if ((priority != nullptr) && (stream != nullptr)) {
-    if (!hip::isValid(stream)) {
-      return hipErrorContextIsDestroyed;
-    }
+    getStreamPerThread(stream);
     *priority = static_cast<int>(reinterpret_cast<hip::Stream*>(stream)->GetPriority());
   } else {
     return hipErrorInvalidValue;