SWDEV-542901-[catch2][dtest]-Stream Capture tests for the hipStreamIsCapturing

projects/hip-tests/catch/unit/graph/hipStreamIsCapturing.cc

@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (c) 2022-2025 Advanced Micro Devices, Inc. All rights reserved.
 Permission is hereby granted, free of charge, to any person obtaining a copy
 of this software and associated documentation files (the "Software"), to deal
 in the Software without restriction, including without limitation the rights
@@ -19,9 +19,9 @@ THE SOFTWARE.
 
 #include <hip_test_common.hh>
 #include <hip_test_kernels.hh>
-
 #include "stream_capture_common.hh"
-
+constexpr size_t N = 100000;
+constexpr size_t Nbytes = N * sizeof(float);
 /**
  * @addtogroup hipStreamIsCapturing hipStreamIsCapturing
  * @{
@@ -53,17 +53,100 @@ TEST_CASE("Unit_hipStreamIsCapturing_Negative_Parameters") {
     HIP_CHECK_ERROR(hipStreamIsCapturing(stream, nullptr), hipErrorInvalidValue);
   }
 
-  SECTION("Check capture status when checked on null stream") {
+  SECTION("Check capture status when checked on different streams") {
     hipStreamCaptureStatus cStatus;
     hipGraph_t graph{nullptr};
-
-    HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal));
-    HIP_CHECK_ERROR(hipStreamIsCapturing(nullptr, &cStatus), hipErrorStreamCaptureImplicit);
+    hipStreamCaptureMode flags = GENERATE(
+        hipStreamCaptureModeGlobal, hipStreamCaptureModeThreadLocal, hipStreamCaptureModeRelaxed);
+    HIP_CHECK(hipStreamBeginCapture(stream, flags));
+    SECTION("Check capture status with null stream") {
+      HIP_CHECK_ERROR(hipStreamIsCapturing(nullptr, &cStatus), hipErrorStreamCaptureImplicit);
+    }
+    SECTION("Check capture status with hipStreamLegacy") {
+      HIP_CHECK_ERROR(hipStreamIsCapturing(hipStreamLegacy, &cStatus),
+                      hipErrorStreamCaptureImplicit);
+    }
     HIP_CHECK(hipStreamEndCapture(stream, &graph));
     HIP_CHECK(hipGraphDestroy(graph));
   }
 }
+/**
+ * Test Description
+ * ------------------------
+ * Enable capture on non-blocking stream
+ * Call  hipStreamGetCaptureInfo with stream nullptr or Legacy Stream.
+ * Returns hipSuccess.
+ * End Capture on stream
+ * Test source
+ * ------------------------
+ * - catch\unit\graph\hipStreamIsCapturing.cc
+ * Test requirements
+ * ------------------------
+ * - HIP_VERSION >= 6.2
+ */
+TEST_CASE("Unit_hipStreamIsCapturing_NegTst_NonBlockingStream") {
+  hipStreamCaptureStatus cStatus;
+  hipStream_t stream;
+  HIP_CHECK(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking));
+  hipGraph_t graph{nullptr};
+  hipStreamCaptureMode flags = GENERATE(hipStreamCaptureModeGlobal, hipStreamCaptureModeThreadLocal,
+                                        hipStreamCaptureModeRelaxed);
+  HIP_CHECK(hipStreamBeginCapture(stream, flags));
+  SECTION("Check capture status with null stream") {
+    HIP_CHECK_ERROR(hipStreamIsCapturing(nullptr, &cStatus), hipSuccess);
+  }
+  SECTION("Check capture status with hipStreamLegacy") {
+    HIP_CHECK_ERROR(hipStreamIsCapturing(hipStreamLegacy, &cStatus), hipSuccess);
+  }
+  HIP_CHECK(hipStreamEndCapture(stream, &graph));
+  HIP_CHECK(hipGraphDestroy(graph));
+  HIP_CHECK(hipStreamIsCapturing(stream, &cStatus));
+  REQUIRE(hipStreamCaptureStatusNone == cStatus);
+}
+/**
+ * Test Description
+ * ------------------------
+ * Enable capture on stream s1
+ * Call APIs that are not supported during capture.
+ * Get capture status must be hiStreamCaptureStatusInvalidated.
+ * End capture on s1
+ * Test source
+ * ------------------------
+ * - catch\unit\graph\hipStreamIsCapturing.cc
+ * Test requirements
+ * ------------------------
+ * - HIP_VERSION >= 6.2
+ */
+TEST_CASE("Unit_hipStreamIsCapturing_NegTst_UnsupportedApi") {
+  hipStream_t stream{nullptr};
+  hipGraph_t graph{nullptr};
+  hipError_t err;
+  HIP_CHECK(hipStreamCreate(&stream));
+  unsigned int* A_mem{nullptr};
+  hipStreamCaptureMode flags = GENERATE(hipStreamCaptureModeGlobal, hipStreamCaptureModeThreadLocal,
+                                        hipStreamCaptureModeRelaxed);
+  HIP_CHECK(hipStreamBeginCapture(stream, flags));
+  hipStreamCaptureStatus captureStatus{hipStreamCaptureStatusNone},
+      captureStatus1{hipStreamCaptureStatusNone};
+  // Check the capture status of the stream
+  HIP_CHECK(hipStreamIsCapturing(stream, &captureStatus));
+  assert(captureStatus == hipStreamCaptureStatusActive);
+  err = hipMalloc(&A_mem, 1024);
+  if (flags == hipStreamCaptureModeRelaxed) {
+    REQUIRE(err == hipSuccess);
+    HIP_CHECK(hipStreamIsCapturing(stream, &captureStatus1));
+    REQUIRE(captureStatus1 == hipStreamCaptureStatusActive);
+    // End the capture
+    HIP_CHECK(hipStreamEndCapture(stream, &graph));
 
+  } else {
+    REQUIRE(err == hipErrorStreamCaptureUnsupported);
+    HIP_CHECK(hipStreamIsCapturing(stream, &captureStatus1));
+    REQUIRE(captureStatus1 == hipStreamCaptureStatusInvalidated);
+    // End the capture
+    HIP_CHECK_ERROR(hipStreamEndCapture(stream, &graph), hipErrorStreamCaptureInvalidated);
+  }
+}
 /**
  * Test Description
  * ------------------------
@@ -197,7 +280,72 @@ TEST_CASE("Unit_hipStreamIsCapturing_Positive_Thread") {
   HIP_CHECK(hipStreamEndCapture(stream, &graph));
   HIP_CHECK(hipGraphDestroy(graph));
 }
-
+/**
+ * Test Description
+ * ------------------------
+ * 1. Create Stream s1 with flag 1 hipStreamDefault (blocking stream)
+ * 2. Enable capture on stream s1.
+ * 3. Enqueue work on null stream should return error.
+ * 4. Query capture status.
+ * 5. End capture on stream s1.
+ * Test source
+ * ------------------------
+ * - catch\unit\graph\hipStreamIsCapturing.cc
+ * Test requirements
+ * ------------------------
+ * - HIP_VERSION >= 6.2
+ */
+#if HT_AMD  // Disabled for NVIDIA due to defect SWDEV-542700
+TEST_CASE("Unit_hipStreamIsCapturing_Unusable_LegacyStream") {
+  hipStream_t stream{nullptr};
+  hipGraph_t graph{nullptr};
+  hipError_t err;
+  HIP_CHECK(hipStreamCreateWithFlags(&stream, hipStreamDefault));
+  float* A_d;
+  float *A_h, *C_h;
+  // Memory allocation to Host pointers
+  A_h = reinterpret_cast<float*>(malloc(Nbytes));
+  C_h = reinterpret_cast<float*>(malloc(Nbytes));
+  REQUIRE(A_h != nullptr);
+  REQUIRE(C_h != nullptr);
+
+  // Memory allocation to Device pointers
+  HIP_CHECK(hipMalloc(&A_d, Nbytes));
+  REQUIRE(A_d != nullptr);
+
+  // Initialize input buffer
+  for (size_t i = 0; i < N; ++i) {
+    A_h[i] = 1.0f + i;
+  }
+  hipStreamCaptureMode flags = GENERATE(hipStreamCaptureModeGlobal, hipStreamCaptureModeThreadLocal,
+                                        hipStreamCaptureModeRelaxed);
+  HIP_CHECK(hipStreamBeginCapture(stream, flags));
+  hipStreamCaptureStatus captureStatus{hipStreamCaptureStatusNone},
+      captureStatus1{hipStreamCaptureStatusNone}, captureStatus2{hipStreamCaptureStatusNone};
+  // Verify the Error returned if null stream is passed.
+  err = hipStreamIsCapturing(0, &captureStatus);
+  REQUIRE(err == hipErrorStreamCaptureImplicit);
+
+  // Check the capture status of the stream
+  HIP_CHECK(hipStreamIsCapturing(stream, &captureStatus1));
+  REQUIRE(captureStatus1 == hipStreamCaptureStatusActive);
+
+  // Copy data to Device
+  HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, 0));
+  HIP_CHECK(hipMemcpyAsync(A_d, C_h, Nbytes, hipMemcpyDeviceToHost, 0));
+
+  // Check the capture status of the stream
+  HIP_CHECK(hipStreamIsCapturing(stream, &captureStatus2));
+  REQUIRE(captureStatus2 == hipStreamCaptureStatusActive);
+
+  // End the capture
+  HIP_CHECK(hipStreamEndCapture(stream, &graph));
+  REQUIRE(graph != nullptr);
+  HIP_CHECK(hipFree(A_d));
+  free(A_h);
+  free(C_h);
+}
+#endif
 /**
  * End doxygen group GraphTest.
  * @}

projects/hip-tests/catch/unit/graph/hipStreamIsCapturing_old.cc

@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (c) 2022-2025 Advanced Micro Devices, Inc. All rights reserved.
 Permission is hereby granted, free of charge, to any person obtaining a copy
 of this software and associated documentation files (the "Software"), to deal
 in the Software without restriction, including without limitation the rights
@@ -89,6 +89,10 @@ TEST_CASE("Unit_hipStreamIsCapturing_Functional_Basic") {
     HIP_CHECK(hipStreamIsCapturing(hipStreamPerThread, &cStatus));
     REQUIRE(hipStreamCaptureStatusNone == cStatus);
   }
+  SECTION("Check capture status with hipStreamLegacy.") {
+    HIP_CHECK(hipStreamIsCapturing(hipStreamLegacy, &cStatus));
+    REQUIRE(hipStreamCaptureStatusNone == cStatus);
+  }
 }
 
 /**