From e6277881c3fcf5b3c1a0c03480bff2c393d742cc Mon Sep 17 00:00:00 2001
From: Vladana Stojiljkovic <Vladana.Stojiljkovic@amd.com>
Date: Tue, 1 Oct 2024 17:19:34 +0200
Subject: [PATCH] SWDEV-486969 - Add macros for capturing sync APIs

Change-Id: I6d348c8b34b29021c281c32aa5a636960e234ccb


[ROCm/hip-tests commit: 70a7d3ab8b356a674c6dc0e65f685bd84e82c330]
---
 .../catch/include/hip_test_common.hh          | 42 +++++++++++++++++++
 1 file changed, 42 insertions(+)

diff --git a/projects/hip-tests/catch/include/hip_test_common.hh b/projects/hip-tests/catch/include/hip_test_common.hh
index 40475f9891..07819ee270 100644
--- a/projects/hip-tests/catch/include/hip_test_common.hh
+++ b/projects/hip-tests/catch/include/hip_test_common.hh
@@ -507,3 +507,45 @@ class BlockingContext {
     HIP_CHECK(hipGraphExecDestroy(graph_exec));                                                    \
     HIP_CHECK(hipGraphDestroy(graph));                                                             \
   }
+
+// These macros are used for testing behaviour when sync APIs are being captured. Before
+// calling BEGIN_CAPTURE_SYNC, hipError_t variable (capture_err) should be initialized to hipSuccess
+// and passed to this macro. The scenario with using this macro should look like this:
+// 1. BEGIN_CAPTURE_SYNC(capture_err)
+// 2. HIP_CHECK_ERROR(SyncAPI, capture_err)
+// 3. END_CAPTURE_SYNC(capture_err)
+// Some sync APIs are allowed in relaxed capture mode which is indicated with
+// rlx_mode_allowed variable. For other two modes, those APIs return
+// hipErrorStreamCaptureUnsupported. These macros shouldn't be used with hipStreamSync and
+// hipDeviceSync during capture.
+#define BEGIN_CAPTURE_SYNC(capture_err, rlx_mode_allowed)                                          \
+  hipStream_t stream;                                                                              \
+  GENERATE_CAPTURE();                                                                              \
+  if (capture) {                                                                                   \
+    HIP_CHECK(hipStreamCreate(&stream));                                                           \
+    hipStreamCaptureMode mode = GENERATE(                                                          \
+        hipStreamCaptureModeGlobal, hipStreamCaptureModeThreadLocal, hipStreamCaptureModeRelaxed); \
+    HIP_CHECK(hipStreamBeginCapture(stream, mode));                                                \
+    if (!rlx_mode_allowed) {                                                                       \
+      capture_err = hipErrorStreamCaptureImplicit;                                                 \
+    } else if (mode != hipStreamCaptureModeRelaxed) {                                              \
+      capture_err = hipErrorStreamCaptureUnsupported;                                              \
+    }                                                                                              \
+  }
+
+// If test has other HIP API calls that depend on sync call that is captured and fails, the rest of
+// the test (except freeing the memory) should be skipped after calling END_CAPTURE_SYNC() by
+// testing if previously created hipError_t variable (capture_err) doesn't equal hipSuccess.
+#define END_CAPTURE_SYNC(capture_err)                                                              \
+  if (capture) {                                                                                   \
+    hipGraph_t graph;                                                                              \
+    hipError_t stream_err = hipSuccess;                                                            \
+    if (capture_err != hipSuccess) {                                                               \
+      stream_err = hipErrorStreamCaptureInvalidated;                                               \
+    }                                                                                              \
+    HIP_CHECK_ERROR(hipStreamEndCapture(stream, &graph), stream_err);                              \
+    if (graph != nullptr) {                                                                        \
+      HIP_CHECK(hipGraphDestroy(graph));                                                           \
+    }                                                                                              \
+    HIP_CHECK(hipStreamDestroy(stream));                                                           \
+  }