From a0e946c2cc55a58bba776c3799640bbe2db2c68a Mon Sep 17 00:00:00 2001 From: "Stojiljkovic, Vladana" Date: Thu, 22 May 2025 17:24:04 +0200 Subject: [PATCH] Match hipStreamAddCallback capture behavior with nvidia (#386) Co-authored-by: Anusha GodavarthySurya --- hipamd/src/hip_stream.cpp | 13 +++++++++++++ 1 file changed, 13 insertions(+) diff --git a/hipamd/src/hip_stream.cpp b/hipamd/src/hip_stream.cpp index a2196cb83a..aec66840b1 100644 --- a/hipamd/src/hip_stream.cpp +++ b/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(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); }