diff --git a/catch/unit/graph/hipStreamIsCapturing.cc b/catch/unit/graph/hipStreamIsCapturing.cc index 4b5abcede..78e31ce7a 100644 --- a/catch/unit/graph/hipStreamIsCapturing.cc +++ b/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 #include - #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,8 +280,73 @@ TEST_CASE("Unit_hipStreamIsCapturing_Positive_Thread") { HIP_CHECK(hipStreamEndCapture(stream, &graph)); HIP_CHECK(hipGraphDestroy(graph)); } - /** -* End doxygen group GraphTest. -* @} -*/ + * 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(malloc(Nbytes)); + C_h = reinterpret_cast(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. + * @} + */ diff --git a/catch/unit/graph/hipStreamIsCapturing_old.cc b/catch/unit/graph/hipStreamIsCapturing_old.cc index d963d43f4..2d1da495c 100644 --- a/catch/unit/graph/hipStreamIsCapturing_old.cc +++ b/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 @@ -56,9 +56,10 @@ Functional Testcase Scenarios : of s1 using hipStreamIsCapturing. Verify that it is in state hipStreamCaptureStatusActive in thread. Exit the thread and end the capture. 11)Functional : Create a stream with default flag (hipStreamDefault). Start capturing - the stream. Invoke hipStreamIsCapturing() on the null stream. Verify hipErrorStreamCaptureImplicit - is returned by hipStreamIsCapturing(). Verify capture status of created stream. Do some operatoins. - End the capture on the created stream. Execute the graph and verify the output from the operations. + the stream. Invoke hipStreamIsCapturing() on the null stream. Verify +hipErrorStreamCaptureImplicit is returned by hipStreamIsCapturing(). Verify capture status of +created stream. Do some operatoins. End the capture on the created stream. Execute the graph and +verify the output from the operations. */ TEST_CASE("Unit_hipStreamIsCapturing_Negative") { @@ -69,8 +70,9 @@ TEST_CASE("Unit_hipStreamIsCapturing_Negative") { ret = hipStreamIsCapturing(stream, nullptr); REQUIRE(hipErrorInvalidValue == ret); } - SECTION("Check capture status with hipStreamPerThread and" - " nullptr as pCaptureStatus.") { + SECTION( + "Check capture status with hipStreamPerThread and" + " nullptr as pCaptureStatus.") { ret = hipStreamIsCapturing(hipStreamPerThread, nullptr); REQUIRE(hipErrorInvalidValue == ret); } @@ -87,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); + } } /** @@ -113,7 +119,7 @@ TEST_CASE("Unit_hipStreamIsCapturing_Functional") { // Fill with Phi + i for (size_t i = 0; i < N; i++) { - A_h[i] = 1.618f + i; + A_h[i] = 1.618f + i; } HIP_CHECK(hipMalloc(&A_d, Nbytes)); @@ -137,8 +143,8 @@ TEST_CASE("Unit_hipStreamIsCapturing_Functional") { HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream)); HIP_CHECK(hipMemsetAsync(C_d, 0, Nbytes, stream)); - hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), - dim3(threadsPerBlock), 0, stream, A_d, C_d, N); + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), dim3(threadsPerBlock), 0, stream, A_d, + C_d, N); HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream)); HIP_CHECK(hipStreamEndCapture(stream, &graph)); @@ -181,7 +187,7 @@ TEST_CASE("Unit_hipStreamIsCapturing_hipStreamPerThread") { // Fill with Phi + i for (size_t i = 0; i < N; i++) { - A_h[i] = 1.618f + i; + A_h[i] = 1.618f + i; } HIP_CHECK(hipMalloc(&A_d, Nbytes)); @@ -194,22 +200,19 @@ TEST_CASE("Unit_hipStreamIsCapturing_hipStreamPerThread") { REQUIRE(hipStreamCaptureStatusNone == cStatus); } - HIP_CHECK(hipStreamBeginCapture(hipStreamPerThread, - hipStreamCaptureModeGlobal)); + HIP_CHECK(hipStreamBeginCapture(hipStreamPerThread, hipStreamCaptureModeGlobal)); SECTION("Start capturing a stream and check the status.") { HIP_CHECK(hipStreamIsCapturing(hipStreamPerThread, &cStatus)); REQUIRE(hipStreamCaptureStatusActive == cStatus); } - HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, - hipStreamPerThread)); + HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, hipStreamPerThread)); HIP_CHECK(hipMemsetAsync(C_d, 0, Nbytes, hipStreamPerThread)); - hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), - dim3(threadsPerBlock), 0, hipStreamPerThread, A_d, C_d, N); - HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, - hipStreamPerThread)); + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), dim3(threadsPerBlock), 0, + hipStreamPerThread, A_d, C_d, N); + HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, hipStreamPerThread)); HIP_CHECK(hipStreamEndCapture(hipStreamPerThread, &graph)); @@ -227,13 +230,13 @@ TEST_CASE("Unit_hipStreamIsCapturing_hipStreamPerThread") { HIP_CHECK(hipFree(C_d)); } /* -* Create 2 streams s1 and s2. Start capturing s1. Record event e1 on s1 and wait -* for event e1 on s2. Queue some operations in s1 and s2. Invoke hipStreamIsCapturing -* on both s1 and s2. Verify that the capture info (status) of both s1 and s2 are identical. -* Record event e2 on s2 and wait for event e2 on s1. End the capture of stream s1. -* Invoke hipStreamIsCapturing on both streams. Verify that the capture info(status) -* of both s1 and s2 are identical. -*/ + * Create 2 streams s1 and s2. Start capturing s1. Record event e1 on s1 and wait + * for event e1 on s2. Queue some operations in s1 and s2. Invoke hipStreamIsCapturing + * on both s1 and s2. Verify that the capture info (status) of both s1 and s2 are identical. + * Record event e2 on s2 and wait for event e2 on s1. End the capture of stream s1. + * Invoke hipStreamIsCapturing on both streams. Verify that the capture info(status) + * of both s1 and s2 are identical. + */ TEST_CASE("Unit_hipStreamIsCapturing_ParentAndForkedStream") { hipStream_t stream1{nullptr}, stream2{nullptr}; hipEvent_t event2{nullptr}, forkStreamEvent{nullptr}; @@ -279,18 +282,17 @@ TEST_CASE("Unit_hipStreamIsCapturing_ParentAndForkedStream") { HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream1)); HIP_CHECK(hipMemcpyAsync(B_d, B_h, Nbytes, hipMemcpyHostToDevice, stream2)); // Kernal Operations - hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), - dim3(threadsPerBlock), 0, stream1, A_d, C_d, N); - hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), - dim3(threadsPerBlock), 0, stream2, B_d, D_d, N); + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), dim3(threadsPerBlock), 0, stream1, A_d, + C_d, N); + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), dim3(threadsPerBlock), 0, stream2, B_d, + D_d, N); // Copy data back to the Host HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream1)); HIP_CHECK(hipMemcpyAsync(D_h, D_d, Nbytes, hipMemcpyDeviceToHost, stream2)); hipStreamCaptureStatus captureStatus1{hipStreamCaptureStatusNone}, - captureStatus2{hipStreamCaptureStatusNone}, - captureStatus3{hipStreamCaptureStatusNone}, - captureStatus4{hipStreamCaptureStatusNone}; + captureStatus2{hipStreamCaptureStatusNone}, captureStatus3{hipStreamCaptureStatusNone}, + captureStatus4{hipStreamCaptureStatusNone}; // Capturing info HIP_CHECK(hipStreamIsCapturing(stream1, &captureStatus1)); HIP_CHECK(hipStreamIsCapturing(stream2, &captureStatus2)); @@ -326,11 +328,11 @@ TEST_CASE("Unit_hipStreamIsCapturing_ParentAndForkedStream") { free(D_h); } /* -* Create a stream s1. Start capturing s1. Get the capture info using hipStreamIsCapturing -* of s1. Launch a thread. In the thread get the capture info of s1 using hipStreamIsCapturing. -* Verify that it is in state hipStreamCaptureStatusActive in thread. Exit the thread and end -* the capture. -*/ + * Create a stream s1. Start capturing s1. Get the capture info using hipStreamIsCapturing + * of s1. Launch a thread. In the thread get the capture info of s1 using hipStreamIsCapturing. + * Verify that it is in state hipStreamCaptureStatusActive in thread. Exit the thread and end + * the capture. + */ // Thread Function static void thread_func(hipStream_t stream) { hipStreamCaptureStatus captureStatus{hipStreamCaptureStatusNone}; @@ -360,11 +362,12 @@ TEST_CASE("Unit_hipStreamIsCapturing_CheckCaptureStatus_FromThread") { } /* -* Create a stream with default flag (hipStreamDefault). Start capturing the stream. -* Invoke hipStreamIsCapturing() on the null stream. Verify hipErrorStreamCaptureImplicit -* is returned by hipStreamIsCapturing(). Verify capture status of created stream. Do some operatoins. -* End the capture on the created stream. Execute the graph and verify the output from the operations. -*/ + * Create a stream with default flag (hipStreamDefault). Start capturing the stream. + * Invoke hipStreamIsCapturing() on the null stream. Verify hipErrorStreamCaptureImplicit + * is returned by hipStreamIsCapturing(). Verify capture status of created stream. Do some + * operatoins. End the capture on the created stream. Execute the graph and verify the output from + * the operations. + */ TEST_CASE("Unit_hipStreamIsCapturing_ChkNullStrmStatus") { hipStream_t stream{nullptr}, streamForGraph{nullptr}; hipGraph_t graph{nullptr}; @@ -394,8 +397,7 @@ TEST_CASE("Unit_hipStreamIsCapturing_ChkNullStrmStatus") { } HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal)); hipStreamCaptureStatus captureStatus{hipStreamCaptureStatusNone}, - captureStatus1{hipStreamCaptureStatusNone}, - captureStatus2{hipStreamCaptureStatusNone}; + captureStatus1{hipStreamCaptureStatusNone}, captureStatus2{hipStreamCaptureStatusNone}; // Verify the Error returned if null stream is passed. ret = hipStreamIsCapturing(0, &captureStatus); REQUIRE(ret == hipErrorStreamCaptureImplicit); @@ -405,8 +407,8 @@ TEST_CASE("Unit_hipStreamIsCapturing_ChkNullStrmStatus") { // Copy data to Device HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream)); // Kernal Operations - hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), - dim3(threadsPerBlock), 0, stream, A_d, C_d, N); + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), dim3(threadsPerBlock), 0, stream, A_d, + C_d, N); HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream)); // End the capture HIP_CHECK(hipStreamEndCapture(stream, &graph)); @@ -423,7 +425,7 @@ TEST_CASE("Unit_hipStreamIsCapturing_ChkNullStrmStatus") { // Verify Output for (size_t i = 0; i < N; i++) { D_h[i] = A_h[i] * A_h[i]; - REQUIRE(C_h[i] == D_h[i]); + REQUIRE(C_h[i] == D_h[i]); } HIP_CHECK(hipGraphExecDestroy(graphExec)); HIP_CHECK(hipGraphDestroy(graph));