From 85f6f2a0c747c1b04a86f1cded05b00b57cb88f7 Mon Sep 17 00:00:00 2001
From: Sarbojit Sarkar <sarbojit.sarkar@amd.com>
Date: Wed, 6 Apr 2022 07:16:18 +0000
Subject: [PATCH] SWDEV-331126 - Fix to Cuda return type

Change-Id: Iea8e7735ef0c58c1244c92338038d91e61d62c16


[ROCm/clr commit: 8657c407212a6f1fc77e4d3164e27b03b11d419c]
---
 projects/clr/hipamd/src/hip_graph.cpp | 17 +++++++++++------
 1 file changed, 11 insertions(+), 6 deletions(-)

diff --git a/projects/clr/hipamd/src/hip_graph.cpp b/projects/clr/hipamd/src/hip_graph.cpp
index 6679c3d506..0e57a6f2ea 100644
--- a/projects/clr/hipamd/src/hip_graph.cpp
+++ b/projects/clr/hipamd/src/hip_graph.cpp
@@ -747,15 +747,20 @@ hipError_t hipStreamBeginCapture(hipStream_t stream, hipStreamCaptureMode mode)
   if (!hip::isValid(stream)) {
     HIP_RETURN(hipErrorInvalidValue);
   }
-  hip::Stream* s = reinterpret_cast<hip::Stream*>(stream);
   // capture cannot be initiated on legacy stream
-  // It can be initiated if the stream is not already in capture mode
-  if (stream == nullptr ||
-      mode < hipStreamCaptureModeGlobal ||
-      mode > hipStreamCaptureModeRelaxed ||
-      s->GetCaptureStatus() == hipStreamCaptureStatusActive) {
+  if (stream == nullptr) {
+    HIP_RETURN(hipErrorStreamCaptureUnsupported);
+  }
+  if (mode < hipStreamCaptureModeGlobal ||
+      mode > hipStreamCaptureModeRelaxed) {
     HIP_RETURN(hipErrorInvalidValue);
   }
+  hip::Stream* s = reinterpret_cast<hip::Stream*>(stream);
+  // It can be initiated if the stream is not already in capture mode
+  if (s->GetCaptureStatus() == hipStreamCaptureStatusActive) {
+    HIP_RETURN(hipErrorIllegalState);
+  }
+
   s->SetCaptureGraph(new ihipGraph());
   s->SetCaptureMode(mode);
   s->SetOriginStream();