From 050d2f991f40627d544b2e12a1dc779b8335c1cd Mon Sep 17 00:00:00 2001 From: Marko Arandjelovic Date: Tue, 1 Oct 2024 17:24:49 +0200 Subject: [PATCH] SWDEV-487395 - Add macros for capturing async APIs Change-Id: Id8647a05f3f483792d9224aaf8a2488c2498c0ba [ROCm/hip-tests commit: aac7b29d7812557f6303a5bd3d34f38a4c9bdc7a] --- .../catch/include/hip_test_common.hh | 23 +++++++++++++++++++ 1 file changed, 23 insertions(+) diff --git a/projects/hip-tests/catch/include/hip_test_common.hh b/projects/hip-tests/catch/include/hip_test_common.hh index 7fec627179..40475f9891 100644 --- a/projects/hip-tests/catch/include/hip_test_common.hh +++ b/projects/hip-tests/catch/include/hip_test_common.hh @@ -484,3 +484,26 @@ class BlockingContext { INFO("Texture is not support on the device. Skipped."); \ return; \ } + +// Call GENERATE_CAPTURE macro at the start of the test, before using BEGIN/END_CAPTURE. +// Use BEGIN/END_CAPTURE macros to execute APIs in both stream capturing and non-capturing modes. +// Place BEGIN_CAPTURE before the API call and END_CAPTURE after the call. +#define GENERATE_CAPTURE() bool capture = GENERATE(true, false); + +#define BEGIN_CAPTURE(stream) \ + if (capture && stream != nullptr) { \ + hipStreamCaptureMode flags = GENERATE( \ + hipStreamCaptureModeGlobal, hipStreamCaptureModeThreadLocal, hipStreamCaptureModeRelaxed); \ + HIP_CHECK(hipStreamBeginCapture(stream, flags)); \ + } + +#define END_CAPTURE(stream) \ + if (capture && stream != nullptr) { \ + hipGraph_t graph = nullptr; \ + hipGraphExec_t graph_exec = nullptr; \ + HIP_CHECK(hipStreamEndCapture(stream, &graph)); \ + HIP_CHECK(hipGraphInstantiate(&graph_exec, graph, nullptr, nullptr, 0)); \ + HIP_CHECK(hipGraphLaunch(graph_exec, stream)); \ + HIP_CHECK(hipGraphExecDestroy(graph_exec)); \ + HIP_CHECK(hipGraphDestroy(graph)); \ + }