Match hipStreamAddCallback capture behavior with nvidia (#386)

Co-authored-by: Anusha GodavarthySurya <Anusha.GodavarthySurya@amd.com>

[ROCm/clr commit: a0e946c2cc]

projects/clr/hipamd/src/hip_stream.cpp

@@ -609,6 +609,19 @@ hipError_t hipStreamAddCallback_common(hipStream_t stream, hipStreamCallback_t c
   if (callback == nullptr || flags != 0) {
     return hipErrorInvalidValue;
   }
+
+  if (stream != nullptr && stream != hipStreamLegacy && hip::isValid(stream)) {
+    hip::Stream* s = reinterpret_cast<hip::Stream*>(stream);
+    if (s->GetCaptureStatus() != hipStreamCaptureStatusNone) {
+      s->SetCaptureStatus(hipStreamCaptureStatusInvalidated);
+      return hipErrorStreamCaptureUnsupported;
+    }
+  } else if (Stream::StreamCaptureBlocking() == true) {
+    // If any of the blocking streams is capturing, return error for implicit capture and
+    // invalidate capture for all capturing streams
+    CHECK_STREAM_CAPTURING();
+  }
+
   StreamCallback* cbo = new StreamAddCallback(stream, callback, userData);
   return streamCallback_common(stream, cbo, userData);
 }