From 16334bc15d58821ad69c2994a741baf588ea0a95 Mon Sep 17 00:00:00 2001
From: "systems-assistant[bot]"
 <221163467+systems-assistant[bot]@users.noreply.github.com>
Date: Wed, 5 Nov 2025 09:32:47 +0100
Subject: [PATCH] SWDEV-525244 - Additional graph validation test (#584)

* SWDEV-525244 - Additional graph validation test

* SWDEV-525244 - Change capture mode to relaxed

* SWDEV-525244 - Additional comments

---------

Co-authored-by: Vladana Stojiljkovic <Vladana.Stojiljkovic@amd.com>
---
 .../catch/unit/graph/hipStreamBeginCapture.cc | 85 +++++++++++++++++++
 1 file changed, 85 insertions(+)

diff --git a/projects/hip-tests/catch/unit/graph/hipStreamBeginCapture.cc b/projects/hip-tests/catch/unit/graph/hipStreamBeginCapture.cc
index f9b8e82121..2721da0504 100644
--- a/projects/hip-tests/catch/unit/graph/hipStreamBeginCapture.cc
+++ b/projects/hip-tests/catch/unit/graph/hipStreamBeginCapture.cc
@@ -1568,6 +1568,91 @@ TEST_CASE("Unit_hipStreamBeginCapture_StreamSync_OngoingCapture_MThread") {
   }
 }
 
+/**
+ * Test Description
+ * ------------------------
+ *    - Test to verify behavior when event is recorded multiple times while capture is active
+ * Test source
+ * ------------------------
+ *    - catch\unit\graph\hipStreamBeginCapture.cc
+ * Test requirements
+ * ------------------------
+ *    - HIP_VERSION >= 5.6
+ */
+TEST_CASE("Unit_hipStreamBeginCapture_MultipleStreams_ReuseEvent") {
+  // Allocate streams
+  hipStream_t str0, str1, str2;
+  hipEvent_t ev0, ev1, ev2;
+  // Instantiate graph for str0
+  hipGraph_t graph0;
+  hipGraphExec_t graphExec0;
+  // Instantiate graph for str1
+  hipGraph_t graph1;
+  hipGraphExec_t graphExec1;
+
+  // Create streams
+  HIP_CHECK(hipStreamCreate(&str0));
+  HIP_CHECK(hipStreamCreate(&str1));
+  HIP_CHECK(hipStreamCreate(&str2));
+
+  // Create events
+  HIP_CHECK(hipEventCreate(&ev0));
+  HIP_CHECK(hipEventCreate(&ev1));
+  HIP_CHECK(hipEventCreate(&ev2));
+
+  // Enable capture on streams str0 and str1
+  HIP_CHECK(hipStreamBeginCapture(str0, hipStreamCaptureModeGlobal));
+  // str1 in relaxed mode so hipGraphInstantiate can be called on cuda without error
+  HIP_CHECK(hipStreamBeginCapture(str1, hipStreamCaptureModeRelaxed));
+
+  dummyKernel<<<1, 1, 0, str0>>>();
+  HIP_CHECK(hipPeekAtLastError());
+  HIP_CHECK(hipEventRecord(ev0, str0));
+  HIP_CHECK(hipEventRecord(ev1, str1));
+  HIP_CHECK(hipStreamWaitEvent(str2, ev0, 0));
+  dummyKernel<<<1, 1, 0, str2>>>();
+  HIP_CHECK(hipPeekAtLastError());
+  HIP_CHECK(hipEventRecord(ev2, str2));
+
+  HIP_CHECK(hipStreamWaitEvent(str0, ev2, 0));
+  dummyKernel<<<1, 1, 0, str0>>>();
+  HIP_CHECK(hipPeekAtLastError());
+
+  // Instantiate graph for str0
+  HIP_CHECK(hipStreamEndCapture(str0, &graph0));
+  HIP_CHECK(hipGraphInstantiate(&graphExec0, graph0, nullptr, nullptr, 0));
+  HIP_CHECK(hipGraphDestroy(graph0));
+
+  HIP_CHECK(hipStreamWaitEvent(str2, ev1, 0));
+  HIP_CHECK(hipEventRecord(ev2, str2));
+
+  HIP_CHECK(hipStreamWaitEvent(str1, ev2, 0));
+  dummyKernel<<<1, 1, 0, str1>>>();
+  HIP_CHECK(hipPeekAtLastError());
+
+  HIP_CHECK(hipStreamEndCapture(str1, &graph1));
+
+  // Launch graph0
+  HIP_CHECK(hipGraphLaunch(graphExec0, str0));
+  HIP_CHECK(hipStreamSynchronize(str0));
+  HIP_CHECK(hipGraphExecDestroy(graphExec0));
+
+  // Instantiate and launch graph for str1
+  HIP_CHECK(hipGraphInstantiate(&graphExec1, graph1, nullptr, nullptr, 0));
+  HIP_CHECK(hipGraphLaunch(graphExec1, str1));
+  HIP_CHECK(hipStreamSynchronize(str1));
+  HIP_CHECK(hipGraphExecDestroy(graphExec1));
+  HIP_CHECK(hipGraphDestroy(graph1));
+
+  // Clean up resources
+  HIP_CHECK(hipEventDestroy(ev0));
+  HIP_CHECK(hipEventDestroy(ev1));
+  HIP_CHECK(hipEventDestroy(ev2));
+  HIP_CHECK(hipStreamDestroy(str0));
+  HIP_CHECK(hipStreamDestroy(str1));
+  HIP_CHECK(hipStreamDestroy(str2));
+}
+
 /**
  * End doxygen group GraphTest.
  * @}