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>
2025-11-05 09:32:47 +01:00
parent 18d4fc460b
commit 16334bc15d
1 changed files with 85 additions and 0 deletions
@@ -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.
 * @}