diff --git a/catch/unit/graph/hipStreamIsCapturing.cc b/catch/unit/graph/hipStreamIsCapturing.cc index 090348a25b..08f59284c8 100644 --- a/catch/unit/graph/hipStreamIsCapturing.cc +++ b/catch/unit/graph/hipStreamIsCapturing.cc @@ -45,9 +45,22 @@ Functional Testcase Scenarios : capture status returned as hipStreamCaptureStatusActive. 8) Functional : Stop capturing using hipStreamPerThread and check status is returned as hipStreamCaptureStatusNone. + 9) Functional : 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 + 10)Functional : 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. + 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. */ - TEST_CASE("Unit_hipStreamIsCapturing_Negative") { hipError_t ret; hipStream_t stream{}; @@ -213,3 +226,211 @@ TEST_CASE("Unit_hipStreamIsCapturing_hipStreamPerThread") { HIP_CHECK(hipFree(A_d)); 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. +*/ +TEST_CASE("Unit_hipStreamIsCapturing_ParentAndForkedStream") { + hipStream_t stream1{nullptr}, stream2{nullptr}; + hipEvent_t event2{nullptr}, forkStreamEvent{nullptr}; + hipGraph_t graph{nullptr}; + constexpr unsigned blocks = 512; + constexpr unsigned threadsPerBlock = 256; + size_t Nbytes = N * sizeof(float); + float *A_d, *B_d, *C_d, *D_d; + float *A_h, *B_h, *C_h, *D_h; + // Memory allocation to Host pointers + A_h = reinterpret_cast(malloc(Nbytes)); + B_h = reinterpret_cast(malloc(Nbytes)); + C_h = reinterpret_cast(malloc(Nbytes)); + D_h = reinterpret_cast(malloc(Nbytes)); + REQUIRE(A_h != nullptr); + REQUIRE(B_h != nullptr); + REQUIRE(C_h != nullptr); + REQUIRE(D_h != nullptr); + // Memory allocation to Device pointers + HIP_CHECK(hipMalloc(&A_d, Nbytes)); + HIP_CHECK(hipMalloc(&B_d, Nbytes)); + HIP_CHECK(hipMalloc(&C_d, Nbytes)); + HIP_CHECK(hipMalloc(&D_d, Nbytes)); + REQUIRE(A_d != nullptr); + REQUIRE(B_d != nullptr); + REQUIRE(C_d != nullptr); + REQUIRE(D_d != nullptr); + + // Initialize input buffer + for (size_t i = 0; i < N; ++i) { + A_h[i] = 3.146f + i; // Pi + B_h[i] = A_h[i]; + } + HIP_CHECK(hipStreamCreate(&stream1)); + HIP_CHECK(hipStreamCreate(&stream2)); + HIP_CHECK(hipEventCreate(&event2)); + HIP_CHECK(hipEventCreate(&forkStreamEvent)); + // Start capture on stream1 + HIP_CHECK(hipStreamBeginCapture(stream1, hipStreamCaptureModeGlobal)); + HIP_CHECK(hipEventRecord(forkStreamEvent, stream1)); + HIP_CHECK(hipStreamWaitEvent(stream2, forkStreamEvent, 0)); + // Copy data to Device + 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); + // 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}; + // Capturing info + HIP_CHECK(hipStreamIsCapturing(stream1, &captureStatus1)); + HIP_CHECK(hipStreamIsCapturing(stream2, &captureStatus2)); + // Verfication of results + REQUIRE(captureStatus1 == hipStreamCaptureStatusActive); + REQUIRE(captureStatus2 == hipStreamCaptureStatusActive); + + HIP_CHECK(hipEventRecord(event2, stream2)); + HIP_CHECK(hipStreamWaitEvent(stream1, event2, 0)); + // End the capture + HIP_CHECK(hipStreamEndCapture(stream1, &graph)); + REQUIRE(graph != nullptr); + + // Capture Info + HIP_CHECK(hipStreamIsCapturing(stream1, &captureStatus3)); + HIP_CHECK(hipStreamIsCapturing(stream2, &captureStatus4)); + // Verification of results + REQUIRE(captureStatus3 == hipStreamCaptureStatusNone); + REQUIRE(captureStatus4 == hipStreamCaptureStatusNone); + + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(stream1)); + HIP_CHECK(hipStreamDestroy(stream2)); + HIP_CHECK(hipEventDestroy(forkStreamEvent)); + HIP_CHECK(hipEventDestroy(event2)); + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipFree(B_d)); + HIP_CHECK(hipFree(C_d)); + HIP_CHECK(hipFree(D_d)); + free(A_h); + free(B_h); + free(C_h); + 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. +*/ +// Thread Function +static void thread_func(hipStream_t stream) { + hipStreamCaptureStatus captureStatus{hipStreamCaptureStatusNone}; + HIP_CHECK(hipStreamIsCapturing(stream, &captureStatus)); + REQUIRE(captureStatus == hipStreamCaptureStatusActive); +} + +TEST_CASE("Unit_hipStreamIsCapturing_CheckCaptureStatus_FromThread") { + hipStream_t stream{nullptr}; + hipGraph_t graph{nullptr}; + + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal)); + // Capture info + hipStreamCaptureStatus captureStatus{hipStreamCaptureStatusNone}; + HIP_CHECK(hipStreamIsCapturing(stream, &captureStatus)); + REQUIRE(captureStatus == hipStreamCaptureStatusActive); + // Thread launch + std::thread t(thread_func, stream); + t.join(); + + HIP_CHECK(hipStreamEndCapture(stream, &graph)); + REQUIRE(graph != nullptr); + + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(stream)); +} + +/* +* 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}; + hipError_t ret; + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipStreamCreate(&streamForGraph)); + float *A_d, *C_d; + float *A_h, *C_h, *D_h; + // Memory allocation to Host pointers + A_h = reinterpret_cast(malloc(Nbytes)); + C_h = reinterpret_cast(malloc(Nbytes)); + D_h = reinterpret_cast(malloc(Nbytes)); + REQUIRE(A_h != nullptr); + REQUIRE(C_h != nullptr); + REQUIRE(D_h != nullptr); + + // Memory allocation to Device pointers + HIP_CHECK(hipMalloc(&A_d, Nbytes)); + HIP_CHECK(hipMalloc(&C_d, Nbytes)); + REQUIRE(A_d != nullptr); + REQUIRE(C_d != nullptr); + + // Initialize input buffer + for (size_t i = 0; i < N; ++i) { + A_h[i] = 1.0f + i; + D_h[i] = 0.0f; + } + HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal)); + hipStreamCaptureStatus captureStatus{hipStreamCaptureStatusNone}, + captureStatus1{hipStreamCaptureStatusNone}, + captureStatus2{hipStreamCaptureStatusNone}; + // Verify the Error returned if null stream is passed. + ret = hipStreamIsCapturing(0, &captureStatus); + REQUIRE(ret == 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, stream)); + // Kernal Operations + 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)); + REQUIRE(graph != nullptr); + + ret = hipStreamIsCapturing(0, &captureStatus2); + REQUIRE(ret == hipSuccess); + + // Launch graph + hipGraphExec_t graphExec; + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + // 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]); + } + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipFree(C_d)); + free(A_h); + free(C_h); + free(D_h); +}