From 76d2c4135e590917d80e039199d6cd7040a49c95 Mon Sep 17 00:00:00 2001 From: "Chaudhary, Jatin Jaikishan" Date: Fri, 1 Aug 2025 10:11:02 +0100 Subject: [PATCH] SWDEV-543350 - restrictions on usage of hipMemcpyWithStream (#727) * SWDEV-543350 - Do not hipMemcpyWithStream to be called during stream capture * update comment --------- Co-authored-by: Jatin Chaudhary --- hipamd/src/hip_memory.cpp | 8 +++++++- 1 file changed, 7 insertions(+), 1 deletion(-) diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index 6dfec59256..280e6252a3 100644 --- a/hipamd/src/hip_memory.cpp +++ b/hipamd/src/hip_memory.cpp @@ -818,11 +818,17 @@ hipError_t hipMemcpy_spt(void* dst, const void* src, size_t sizeBytes, hipMemcpy hipError_t hipMemcpyWithStream(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, hipStream_t stream) { HIP_INIT_API(hipMemcpyWithStream, dst, src, sizeBytes, kind, stream); - STREAM_CAPTURE(hipMemcpyAsync, stream, dst, src, sizeBytes, kind); if (!hip::isValid(stream)) { HIP_RETURN(hipErrorContextIsDestroyed); } + // Since this API does a MemcpyAsync + sync, and sync during a graph capture + // is not allowed, we raise an error here if user tries to use this during a + // stream capture. + if (hip::Stream::StreamCaptureOngoing(stream) == true) { + HIP_RETURN(hipErrorStreamCaptureUnsupported); + } + hip::Stream* hip_stream = hip::getStream(stream); if (hip_stream == nullptr) { HIP_RETURN(hipErrorInvalidValue);