From 331a2815a7c91852e010e35e2e360c2febab1a14 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Wed, 21 Jun 2023 15:59:45 +0530 Subject: [PATCH] SWDEV-389399 - [catch2][dtest] Adding hipStreamSynchronize scenarios when streams are in capturing mode (#270) Change-Id: Ib2f919d84f6acc136db582ca0ac717634e8e5be9 [ROCm/hip-tests commit: b65b09234dc54e9b1bfc2320bcea81d81805e4e2] --- .../catch/unit/graph/hipStreamBeginCapture.cc | 449 ++++++++++++++---- 1 file changed, 361 insertions(+), 88 deletions(-) diff --git a/projects/hip-tests/catch/unit/graph/hipStreamBeginCapture.cc b/projects/hip-tests/catch/unit/graph/hipStreamBeginCapture.cc index 81f97e5f92..cc7e85ef5f 100644 --- a/projects/hip-tests/catch/unit/graph/hipStreamBeginCapture.cc +++ b/projects/hip-tests/catch/unit/graph/hipStreamBeginCapture.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2022 - 2023 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 @@ -20,8 +20,7 @@ THE SOFTWARE. #include #include #include - -#include "stream_capture_common.hh" +#include "stream_capture_common.hh" // NOLINT /** * @addtogroup hipStreamBeginCapture hipStreamBeginCapture @@ -56,7 +55,8 @@ static void hostNodeCallback(void* data) { } template -void captureStreamAndLaunchGraph(F graphFunc, hipStreamCaptureMode mode, hipStream_t stream) { +void captureStreamAndLaunchGraph(F graphFunc, hipStreamCaptureMode mode, + hipStream_t stream) { constexpr size_t N = 1000000; size_t Nbytes = N * sizeof(T); @@ -88,7 +88,8 @@ void captureStreamAndLaunchGraph(F graphFunc, hipStreamCaptureMode mode, hipStre std::fill_n(A_h.host_ptr(), N, static_cast(i)); HIP_CHECK(hipGraphLaunch(graphExec, stream)); HIP_CHECK(hipStreamSynchronize(stream)); - ArrayFindIfNot(B_h.host_ptr(), static_cast(i) * static_cast(i), N); + ArrayFindIfNot(B_h.host_ptr(), + static_cast(i) * static_cast(i), N); } HIP_CHECK(hipGraphExecDestroy(graphExec)) @@ -115,15 +116,16 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_Functional") { StreamGuard stream_guard(stream_type); hipStream_t stream = stream_guard.stream(); - const hipStreamCaptureMode captureMode = GENERATE( - hipStreamCaptureModeGlobal, hipStreamCaptureModeThreadLocal, hipStreamCaptureModeRelaxed); + const hipStreamCaptureMode captureMode = GENERATE(hipStreamCaptureModeGlobal, + hipStreamCaptureModeThreadLocal, hipStreamCaptureModeRelaxed); EventsGuard events_guard(3); StreamsGuard streams_guard(2); SECTION("Linear graph capture") { captureStreamAndLaunchGraph( - [](float* A_h, float* A_d, float* B_h, float* B_d, size_t N, hipStream_t stream) { + [](float* A_h, float* A_d, float* B_h, float* B_d, size_t N, + hipStream_t stream) { return captureSequenceLinear(A_h, A_d, B_h, B_d, N, stream); }, captureMode, stream); @@ -131,10 +133,10 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_Functional") { SECTION("Branched graph capture") { captureStreamAndLaunchGraph( - [&streams_guard, &events_guard](float* A_h, float* A_d, float* B_h, float* B_d, size_t N, - hipStream_t stream) { - captureSequenceBranched(A_h, A_d, B_h, B_d, N, stream, streams_guard.stream_list(), - events_guard.event_list()); + [&streams_guard, &events_guard](float* A_h, float* A_d, float* B_h, + float* B_d, size_t N, hipStream_t stream) { + captureSequenceBranched(A_h, A_d, B_h, B_d, N, stream, + streams_guard.stream_list(), events_guard.event_list()); }, captureMode, stream); } @@ -170,7 +172,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_Negative_Parameters") { hipErrorIllegalState); } SECTION("Creating hipStream with invalid mode") { - HIP_CHECK_ERROR(hipStreamBeginCapture(stream, hipStreamCaptureMode(-1)), hipErrorInvalidValue); + HIP_CHECK_ERROR(hipStreamBeginCapture(stream, hipStreamCaptureMode(-1)), + hipErrorInvalidValue); } #if HT_NVIDIA // EXSWHTEC-216 SECTION("Stream capture on uninitialized stream returns error code.") { @@ -178,7 +181,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_Negative_Parameters") { StreamGuard sg(Streams::created); return sg.stream(); }; - HIP_CHECK_ERROR(hipStreamBeginCapture(InvalidStream(), hipStreamCaptureModeGlobal), + HIP_CHECK_ERROR(hipStreamBeginCapture(InvalidStream(), + hipStreamCaptureModeGlobal), hipErrorContextIsDestroyed); } #endif @@ -202,8 +206,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_Basic") { StreamGuard stream_guard(stream_type); hipStream_t s = stream_guard.stream(); - const hipStreamCaptureMode captureMode = GENERATE( - hipStreamCaptureModeGlobal, hipStreamCaptureModeThreadLocal, hipStreamCaptureModeRelaxed); + const hipStreamCaptureMode captureMode = GENERATE(hipStreamCaptureModeGlobal, + hipStreamCaptureModeThreadLocal, hipStreamCaptureModeRelaxed); HIP_CHECK(hipStreamBeginCapture(s, captureMode)); @@ -213,7 +217,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_Basic") { /* Local function for inter stream event synchronization */ -static void interStrmEventSyncCapture(const hipStream_t& stream1, const hipStream_t& stream2) { +static void interStrmEventSyncCapture(const hipStream_t& stream1, + const hipStream_t& stream2) { hipGraph_t graph1{nullptr}, graph2{nullptr}; hipGraphExec_t graphExec1{nullptr}, graphExec2{nullptr}; @@ -260,7 +265,8 @@ static void interStrmEventSyncCapture(const hipStream_t& stream1, const hipStrea /* Local function for colligated stream capture */ -static void colligatedStrmCapture(const hipStream_t& stream1, const hipStream_t& stream2) { +static void colligatedStrmCapture(const hipStream_t& stream1, + const hipStream_t& stream2) { hipGraph_t graph1{nullptr}, graph2{nullptr}; hipGraphExec_t graphExec1{nullptr}, graphExec2{nullptr}; @@ -303,7 +309,8 @@ static void colligatedStrmCapture(const hipStream_t& stream1, const hipStream_t& /* Local function for colligated stream capture functionality */ -static void colligatedStrmCaptureFunc(const hipStream_t& stream1, const hipStream_t& stream2) { +static void colligatedStrmCaptureFunc(const hipStream_t& stream1, + const hipStream_t& stream2) { constexpr size_t N = 1000000; size_t Nbytes = N * sizeof(int); @@ -323,8 +330,10 @@ static void colligatedStrmCaptureFunc(const hipStream_t& stream1, const hipStrea // Capture 2 streams HIP_CHECK(hipStreamBeginCapture(stream1, hipStreamCaptureModeGlobal)); HIP_CHECK(hipStreamBeginCapture(stream2, hipStreamCaptureModeGlobal)); - captureSequenceLinear(A_h.host_ptr(), A_d.ptr(), B_h.host_ptr(), B_d.ptr(), N, stream1); - captureSequenceLinear(C_h.host_ptr(), C_d.ptr(), D_h.host_ptr(), D_d.ptr(), N, stream2); + captureSequenceLinear(A_h.host_ptr(), A_d.ptr(), B_h.host_ptr(), B_d.ptr(), + N, stream1); + captureSequenceLinear(C_h.host_ptr(), C_d.ptr(), D_h.host_ptr(), D_d.ptr(), + N, stream2); captureSequenceCompute(A_d.ptr(), B_h.host_ptr(), B_d.ptr(), N, stream1); captureSequenceCompute(C_d.ptr(), D_h.host_ptr(), D_d.ptr(), N, stream2); HIP_CHECK(hipStreamEndCapture(stream1, &graph1)); @@ -360,8 +369,9 @@ static void colligatedStrmCaptureFunc(const hipStream_t& stream1, const hipStrea /* Stream Capture thread function */ -static void threadStrmCaptureFunc(hipStream_t stream, int* A_h, int* A_d, int* B_h, int* B_d, - hipGraph_t* graph, size_t N, hipStreamCaptureMode mode) { +static void threadStrmCaptureFunc(hipStream_t stream, int* A_h, int* A_d, + int* B_h, int* B_d, hipGraph_t* graph, + size_t N, hipStreamCaptureMode mode) { // Capture stream HIP_CHECK(hipStreamBeginCapture(stream, mode)); captureSequenceLinear(A_h, A_d, B_h, B_d, N, stream); @@ -393,10 +403,10 @@ static void multithreadedTest(hipStreamCaptureMode mode) { LinearAllocGuard D_d(LinearAllocs::hipMalloc, Nbytes); // Launch 2 threads to capture the 2 streams into graphs - std::thread t1(threadStrmCaptureFunc, stream1, A_h.host_ptr(), A_d.ptr(), B_h.host_ptr(), - B_d.ptr(), &graph1, N, mode); - std::thread t2(threadStrmCaptureFunc, stream2, C_h.host_ptr(), C_d.ptr(), D_h.host_ptr(), - D_d.ptr(), &graph2, N, mode); + std::thread t1(threadStrmCaptureFunc, stream1, A_h.host_ptr(), A_d.ptr(), + B_h.host_ptr(), B_d.ptr(), &graph1, N, mode); + std::thread t2(threadStrmCaptureFunc, stream2, C_h.host_ptr(), C_d.ptr(), + D_h.host_ptr(), D_d.ptr(), &graph2, N, mode); t1.join(); t2.join(); @@ -469,9 +479,11 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_InterStrmEventSync_Flags") { TEST_CASE("Unit_hipStreamBeginCapture_Positive_InterStrmEventSync_Priority") { int minPriority = 0, maxPriority = 0; HIP_CHECK(hipDeviceGetStreamPriorityRange(&minPriority, &maxPriority)); - StreamGuard stream_guard1(Streams::withPriority, hipStreamDefault, minPriority); + StreamGuard stream_guard1(Streams::withPriority, hipStreamDefault, + minPriority); hipStream_t stream1 = stream_guard1.stream(); - StreamGuard stream_guard2(Streams::withPriority, hipStreamDefault, maxPriority); + StreamGuard stream_guard2(Streams::withPriority, hipStreamDefault, + maxPriority); hipStream_t stream2 = stream_guard2.stream(); interStrmEventSyncCapture(stream1, stream2); } @@ -517,12 +529,14 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_ColligatedStrmCapture_Flags") { * ------------------------ * - HIP_VERSION >= 5.2 */ -TEST_CASE("Unit_hipStreamBeginCapture_Positive_ColligatedStrmCapture_Priority") { +TEST_CASE("Unit_hipStreamBeginCapture_Positive_ColligatedStrmCapture_Prio") { int minPriority = 0, maxPriority = 0; HIP_CHECK(hipDeviceGetStreamPriorityRange(&minPriority, &maxPriority)); - StreamGuard stream_guard1(Streams::withPriority, hipStreamDefault, minPriority); + StreamGuard stream_guard1(Streams::withPriority, hipStreamDefault, + minPriority); hipStream_t stream1 = stream_guard1.stream(); - StreamGuard stream_guard2(Streams::withPriority, hipStreamDefault, maxPriority); + StreamGuard stream_guard2(Streams::withPriority, hipStreamDefault, + maxPriority); hipStream_t stream2 = stream_guard2.stream(); colligatedStrmCapture(stream1, stream2); } @@ -563,8 +577,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_ColligatedStrmCaptureFunc") { * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_hipStreamBeginCapture_Positive_Multithreaded") { - const hipStreamCaptureMode captureMode = GENERATE( - hipStreamCaptureModeGlobal, hipStreamCaptureModeThreadLocal, hipStreamCaptureModeRelaxed); + const hipStreamCaptureMode captureMode = GENERATE(hipStreamCaptureModeGlobal, + hipStreamCaptureModeThreadLocal, hipStreamCaptureModeRelaxed); multithreadedTest(captureMode); } @@ -693,7 +707,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_CapturingFromWithinStrms") { HIP_CHECK(hipEventRecord(events[2], streams[2])); HIP_CHECK(hipStreamWaitEvent(streams[0], events[1], 0)); HIP_CHECK(hipStreamWaitEvent(streams[0], events[2], 0)); - HIP_CHECK(hipMemcpyAsync(hostMem, devMem, sizeof(int), hipMemcpyDefault, streams[0])); + HIP_CHECK(hipMemcpyAsync(hostMem, devMem, sizeof(int), hipMemcpyDefault, + streams[0])); HIP_CHECK(hipStreamEndCapture(streams[0], &graph)); // End Capture // Reset device memory HIP_CHECK(hipMemset(devMem, 0, sizeof(int))); @@ -735,7 +750,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_Negative_DetectingInvalidCapture") { dummyKernel<<<1, 1, 0, streams[0]>>>(); // Since stream[1] is already in capture mode due to event wait // hipStreamBeginCapture on stream[1] is expected to return error. - HIP_CHECK_ERROR(hipStreamBeginCapture(streams[1], hipStreamCaptureModeGlobal), + HIP_CHECK_ERROR(hipStreamBeginCapture(streams[1], + hipStreamCaptureModeGlobal), hipErrorIllegalState); } @@ -768,7 +784,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_CapturingMultGraphsFrom1Strm") { for (int i = 0; i < 3; i++) { HIP_CHECK(hipStreamBeginCapture(stream1, hipStreamCaptureModeGlobal)); for (int j = 0; j <= i; j++) incrementKernel<<<1, 1, 0, stream1>>>(devMem); - HIP_CHECK(hipMemcpyAsync(hostMem, devMem, sizeof(int), hipMemcpyDefault, stream1)); + HIP_CHECK(hipMemcpyAsync(hostMem, devMem, sizeof(int), hipMemcpyDefault, + stream1)); HIP_CHECK(hipStreamEndCapture(stream1, &graphs[i])); } // Instantiate and execute all graphs @@ -784,7 +801,6 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_CapturingMultGraphsFrom1Strm") { } } -#if HT_NVIDIA /** * Test Description * ------------------------ @@ -808,29 +824,35 @@ TEST_CASE("Unit_hipStreamBeginCapture_Negative_CheckingSyncDuringCapture") { EventsGuard events_guard(1); hipEvent_t e = events_guard[0]; - const hipStreamCaptureMode captureMode = GENERATE( - hipStreamCaptureModeGlobal, hipStreamCaptureModeThreadLocal, hipStreamCaptureModeRelaxed); + const hipStreamCaptureMode captureMode = GENERATE(hipStreamCaptureModeGlobal, + hipStreamCaptureModeThreadLocal, hipStreamCaptureModeRelaxed); HIP_CHECK(hipStreamBeginCapture(stream, captureMode)); SECTION("Synchronize stream during capture") { - HIP_CHECK_ERROR(hipStreamSynchronize(stream), hipErrorStreamCaptureUnsupported); + HIP_CHECK_ERROR(hipStreamSynchronize(stream), + hipErrorStreamCaptureUnsupported); } + SECTION("Query stream during capture") { + HIP_CHECK_ERROR(hipStreamQuery(stream), + hipErrorStreamCaptureUnsupported); + } +#if HT_NVIDIA SECTION("Synchronize device during capture") { - HIP_CHECK_ERROR(hipDeviceSynchronize(), hipErrorStreamCaptureUnsupported); + HIP_CHECK_ERROR(hipDeviceSynchronize(), + hipErrorStreamCaptureUnsupported); } SECTION("Synchronize event during capture") { HIP_CHECK(hipEventRecord(e, stream)); HIP_CHECK_ERROR(hipEventSynchronize(e), hipErrorCapturedEvent); } - SECTION("Query stream during capture") { - HIP_CHECK_ERROR(hipStreamQuery(stream), hipErrorStreamCaptureUnsupported); - } SECTION("Query for an event during capture") { HIP_CHECK(hipEventRecord(e, stream)); HIP_CHECK_ERROR(hipEventQuery(e), hipErrorCapturedEvent); } +#endif } +#if HT_NVIDIA /** * Test Description * ------------------------ @@ -861,14 +883,17 @@ TEST_CASE("Unit_hipStreamBeginCapture_Negative_UnsafeCallsDuringCapture") { HIP_CHECK(hipStreamBeginCapture(stream, captureMode)); SECTION("hipMalloc during capture") { - HIP_CHECK_ERROR(hipMalloc(&devMem2, sizeof(int)), hipErrorStreamCaptureUnsupported); + HIP_CHECK_ERROR(hipMalloc(&devMem2, sizeof(int)), + hipErrorStreamCaptureUnsupported); } SECTION("hipMemcpy during capture") { - HIP_CHECK_ERROR(hipMemcpy(devMem.ptr(), hostMem.host_ptr(), sizeof(int), hipMemcpyHostToDevice), + HIP_CHECK_ERROR(hipMemcpy(devMem.ptr(), hostMem.host_ptr(), sizeof(int), + hipMemcpyHostToDevice), hipErrorStreamCaptureImplicit); } SECTION("hipMemset during capture") { - HIP_CHECK_ERROR(hipMemset(devMem.ptr(), 0, sizeof(int)), hipErrorStreamCaptureImplicit); + HIP_CHECK_ERROR(hipMemset(devMem.ptr(), 0, sizeof(int)), + hipErrorStreamCaptureImplicit); } } #endif @@ -889,7 +914,7 @@ TEST_CASE("Unit_hipStreamBeginCapture_Negative_UnsafeCallsDuringCapture") { * ------------------------ * - HIP_VERSION >= 5.2 */ -TEST_CASE("Unit_hipStreamBeginCapture_Negative_EndingCapturewhenCaptureInProgress") { +TEST_CASE("Unit_hipStreamBeginCapture_Negative_EndingCapwhenCapInProg") { hipGraph_t graph{nullptr}; StreamsGuard streams_guard(2); @@ -905,7 +930,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_Negative_EndingCapturewhenCaptureInProgres HIP_CHECK(hipEventRecord(e, stream1)); HIP_CHECK(hipStreamWaitEvent(stream2, e, 0)); dummyKernel<<<1, 1, 0, stream2>>>(); - HIP_CHECK_ERROR(hipStreamEndCapture(stream1, &graph), hipErrorStreamCaptureUnjoined); + HIP_CHECK_ERROR(hipStreamEndCapture(stream1, &graph), + hipErrorStreamCaptureUnjoined); } SECTION("End strm capture when forked strm still has operations") { EventsGuard events_guard(2); @@ -919,7 +945,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_Negative_EndingCapturewhenCaptureInProgres HIP_CHECK(hipEventRecord(e2, stream2)); HIP_CHECK(hipStreamWaitEvent(stream1, e2, 0)); dummyKernel<<<1, 1, 0, stream2>>>(); - HIP_CHECK_ERROR(hipStreamEndCapture(stream1, &graph), hipErrorStreamCaptureUnjoined); + HIP_CHECK_ERROR(hipStreamEndCapture(stream1, &graph), + hipErrorStreamCaptureUnjoined); } } /** @@ -942,17 +969,19 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_MultiGPU") { SUCCEED("skipping the testcases as numDevices < 2"); return; } - hipStream_t* stream = reinterpret_cast(malloc(devcount * sizeof(hipStream_t))); + hipStream_t* stream = reinterpret_cast + (malloc(devcount * sizeof(hipStream_t))); REQUIRE(stream != nullptr); - hipGraph_t* graph = reinterpret_cast(malloc(devcount * sizeof(hipGraph_t))); + hipGraph_t* graph = reinterpret_cast + (malloc(devcount * sizeof(hipGraph_t))); REQUIRE(graph != nullptr); int **devMem{nullptr}, **hostMem{nullptr}; hostMem = reinterpret_cast(malloc(sizeof(int*) * devcount)); REQUIRE(hostMem != nullptr); devMem = reinterpret_cast(malloc(sizeof(int*) * devcount)); REQUIRE(devMem != nullptr); - hipGraphExec_t* graphExec = - reinterpret_cast(malloc(devcount * sizeof(hipGraphExec_t))); + hipGraphExec_t* graphExec = reinterpret_cast + (malloc(devcount * sizeof(hipGraphExec_t))); // Capture stream in each device for (int dev = 0; dev < devcount; dev++) { HIP_CHECK(hipSetDevice(dev)); @@ -964,14 +993,15 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_MultiGPU") { for (int i = 0; i < (dev + 1); i++) { incrementKernel<<<1, 1, 0, stream[dev]>>>(devMem[dev]); } - HIP_CHECK( - hipMemcpyAsync(hostMem[dev], devMem[dev], sizeof(int), hipMemcpyDefault, stream[dev])); + HIP_CHECK(hipMemcpyAsync(hostMem[dev], devMem[dev], sizeof(int), + hipMemcpyDefault, stream[dev])); HIP_CHECK(hipStreamEndCapture(stream[dev], &graph[dev])); } // Launch the captured graphs in the respective device for (int dev = 0; dev < devcount; dev++) { HIP_CHECK(hipSetDevice(dev)); - HIP_CHECK(hipGraphInstantiate(&graphExec[dev], graph[dev], nullptr, nullptr, 0)); + HIP_CHECK(hipGraphInstantiate(&graphExec[dev], graph[dev], nullptr, + nullptr, 0)); HIP_CHECK(hipGraphLaunch(graphExec[dev], stream[dev])); } // Validate output @@ -1038,8 +1068,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_nestedStreamCapture") { HIP_CHECK(hipEventRecord(events[3], streams[2])); HIP_CHECK(hipStreamWaitEvent(streams[0], events[3], 0)); HIP_CHECK(hipStreamWaitEvent(streams[0], events[2], 0)); - HIP_CHECK(hipMemcpyAsync(hostMem_g.host_ptr(), devMem_g.ptr(), sizeof(int), hipMemcpyDefault, - streams[0])); + HIP_CHECK(hipMemcpyAsync(hostMem_g.host_ptr(), devMem_g.ptr(), sizeof(int), + hipMemcpyDefault, streams[0])); HIP_CHECK(hipStreamEndCapture(streams[0], &graph)); // End Capture // Reset device memory HIP_CHECK(hipMemset(devMem_g.ptr(), 0, sizeof(int))); @@ -1077,15 +1107,23 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_streamReuse") { hipGraph_t graphs[3]; StreamsGuard streams(3); EventsGuard events(4); - LinearAllocGuard hostMem_g1 = LinearAllocGuard(LinearAllocs::malloc, sizeof(int)); - LinearAllocGuard hostMem_g2 = LinearAllocGuard(LinearAllocs::malloc, sizeof(int)); - LinearAllocGuard hostMem_g3 = LinearAllocGuard(LinearAllocs::malloc, sizeof(int)); - LinearAllocGuard devMem_g1 = LinearAllocGuard(LinearAllocs::hipMalloc, sizeof(int)); - LinearAllocGuard devMem_g2 = LinearAllocGuard(LinearAllocs::hipMalloc, sizeof(int)); - LinearAllocGuard devMem_g3 = LinearAllocGuard(LinearAllocs::hipMalloc, sizeof(int)); + LinearAllocGuard hostMem_g1 = LinearAllocGuard + (LinearAllocs::malloc, sizeof(int)); + LinearAllocGuard hostMem_g2 = LinearAllocGuard + (LinearAllocs::malloc, sizeof(int)); + LinearAllocGuard hostMem_g3 = LinearAllocGuard + (LinearAllocs::malloc, sizeof(int)); + LinearAllocGuard devMem_g1 = LinearAllocGuard + (LinearAllocs::hipMalloc, sizeof(int)); + LinearAllocGuard devMem_g2 = LinearAllocGuard + (LinearAllocs::hipMalloc, sizeof(int)); + LinearAllocGuard devMem_g3 = LinearAllocGuard + (LinearAllocs::hipMalloc, sizeof(int)); - std::vector hostMem = {hostMem_g1.host_ptr(), hostMem_g2.host_ptr(), hostMem_g3.host_ptr()}; - std::vector devMem = {devMem_g1.ptr(), devMem_g2.ptr(), devMem_g3.ptr()}; + std::vector hostMem = {hostMem_g1.host_ptr(), hostMem_g2.host_ptr(), + hostMem_g3.host_ptr()}; + std::vector devMem = {devMem_g1.ptr(), devMem_g2.ptr(), + devMem_g3.ptr()}; // Create a device memory of size int and initialize it to 0 for (int i = 0; i < 3; i++) { memset(hostMem[i], 0, sizeof(int)); @@ -1109,14 +1147,16 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_streamReuse") { HIP_CHECK(hipEventRecord(events[3], streams[2])); HIP_CHECK(hipStreamWaitEvent(streams[0], events[3], 0)); HIP_CHECK(hipStreamWaitEvent(streams[0], events[2], 0)); - HIP_CHECK(hipMemcpyAsync(hostMem[0], devMem[0], sizeof(int), hipMemcpyDefault, streams[0])); + HIP_CHECK(hipMemcpyAsync(hostMem[0], devMem[0], sizeof(int), + hipMemcpyDefault, streams[0])); HIP_CHECK(hipStreamEndCapture(streams[0], &graphs[0])); // End Capture // Start capturing graph2 from stream 2 HIP_CHECK(hipStreamBeginCapture(streams[1], hipStreamCaptureModeGlobal)); incrementKernel<<<1, 1, 0, streams[1]>>>(devMem[1]); incrementKernel<<<1, 1, 0, streams[1]>>>(devMem[1]); incrementKernel<<<1, 1, 0, streams[1]>>>(devMem[1]); - HIP_CHECK(hipMemcpyAsync(hostMem[1], devMem[1], sizeof(int), hipMemcpyDefault, streams[1])); + HIP_CHECK(hipMemcpyAsync(hostMem[1], devMem[1], sizeof(int), + hipMemcpyDefault, streams[1])); HIP_CHECK(hipStreamEndCapture(streams[1], &graphs[1])); // End Capture // Start capturing graph3 from stream 3 HIP_CHECK(hipStreamBeginCapture(streams[2], hipStreamCaptureModeGlobal)); @@ -1125,7 +1165,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_streamReuse") { incrementKernel<<<1, 1, 0, streams[2]>>>(devMem[2]); incrementKernel<<<1, 1, 0, streams[2]>>>(devMem[2]); incrementKernel<<<1, 1, 0, streams[2]>>>(devMem[2]); - HIP_CHECK(hipMemcpyAsync(hostMem[2], devMem[2], sizeof(int), hipMemcpyDefault, streams[2])); + HIP_CHECK(hipMemcpyAsync(hostMem[2], devMem[2], sizeof(int), + hipMemcpyDefault, streams[2])); HIP_CHECK(hipStreamEndCapture(streams[2], &graphs[2])); // End Capture // Reset device memory HIP_CHECK(hipMemset(devMem[0], 0, sizeof(int))); @@ -1169,32 +1210,40 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_captureComplexGraph") { EventsGuard events(7); // Allocate Device memory and Host memory size_t N = GRIDSIZE * BLOCKSIZE; - LinearAllocGuard Ah = LinearAllocGuard(LinearAllocs::malloc, N * sizeof(int)); - LinearAllocGuard Bh = LinearAllocGuard(LinearAllocs::malloc, N * sizeof(int)); - LinearAllocGuard Ch = LinearAllocGuard(LinearAllocs::malloc, N * sizeof(int)); - LinearAllocGuard Ad = LinearAllocGuard(LinearAllocs::hipMalloc, N * sizeof(int)); - LinearAllocGuard Bd = LinearAllocGuard(LinearAllocs::hipMalloc, N * sizeof(int)); + LinearAllocGuard Ah = LinearAllocGuard + (LinearAllocs::malloc, N * sizeof(int)); + LinearAllocGuard Bh = LinearAllocGuard + (LinearAllocs::malloc, N * sizeof(int)); + LinearAllocGuard Ch = LinearAllocGuard + (LinearAllocs::malloc, N * sizeof(int)); + LinearAllocGuard Ad = LinearAllocGuard + (LinearAllocs::hipMalloc, N * sizeof(int)); + LinearAllocGuard Bd = LinearAllocGuard + (LinearAllocs::hipMalloc, N * sizeof(int)); // Capture streams into graph HIP_CHECK(hipStreamBeginCapture(streams[0], hipStreamCaptureModeGlobal)); HIP_CHECK(hipEventRecord(events[0], streams[0])); HIP_CHECK(hipStreamWaitEvent(streams[3], events[0], 0)); HIP_CHECK(hipStreamWaitEvent(streams[4], events[0], 0)); - HIP_CHECK( - hipMemcpyAsync(Ad.ptr(), Ah.host_ptr(), (N * sizeof(int)), hipMemcpyDefault, streams[0])); - HIP_CHECK( - hipMemcpyAsync(Bd.ptr(), Bh.host_ptr(), (N * sizeof(int)), hipMemcpyDefault, streams[4])); + HIP_CHECK(hipMemcpyAsync(Ad.ptr(), Ah.host_ptr(), (N * sizeof(int)), + hipMemcpyDefault, streams[0])); + HIP_CHECK(hipMemcpyAsync(Bd.ptr(), Bh.host_ptr(), (N * sizeof(int)), + hipMemcpyDefault, streams[4])); hipHostFn_t fn = hostNodeCallback; HIPCHECK(hipLaunchHostFunc(streams[3], fn, nullptr)); HIP_CHECK(hipEventRecord(events[1], streams[0])); HIP_CHECK(hipStreamWaitEvent(streams[1], events[1], 0)); int* Ad_2nd_half = Ad.ptr() + N / 2; int* Ad_1st_half = Ad.ptr(); - mymul<<>>(Ad_2nd_half, CONST_KER2_VAL); - mymul<<>>(Ad_1st_half, CONST_KER1_VAL); + mymul<<>>(Ad_2nd_half, + CONST_KER2_VAL); + mymul<<>>(Ad_1st_half, + CONST_KER1_VAL); HIP_CHECK(hipEventRecord(events[2], streams[1])); HIP_CHECK(hipStreamWaitEvent(streams[2], events[2], 0)); - mymul<<>>(Ad_1st_half, CONST_KER3_VAL); + mymul<<>>(Ad_1st_half, + CONST_KER3_VAL); HIPCHECK(hipLaunchHostFunc(streams[2], fn, nullptr)); HIP_CHECK(hipEventRecord(events[6], streams[1])); HIP_CHECK(hipStreamWaitEvent(streams[0], events[6], 0)); @@ -1205,8 +1254,8 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_captureComplexGraph") { HIP_CHECK(hipStreamWaitEvent(streams[0], events[3], 0)); HIP_CHECK(hipEventRecord(events[4], streams[3])); HIP_CHECK(hipStreamWaitEvent(streams[0], events[4], 0)); - HIP_CHECK( - hipMemcpyAsync(Ch.host_ptr(), Ad.ptr(), (N * sizeof(int)), hipMemcpyDefault, streams[0])); + HIP_CHECK(hipMemcpyAsync(Ch.host_ptr(), Ad.ptr(), (N * sizeof(int)), + hipMemcpyDefault, streams[0])); HIP_CHECK(hipStreamEndCapture(streams[0], &graph)); // End Capture // Execute and test the graph hipGraphExec_t graphExec{nullptr}; @@ -1219,10 +1268,11 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_captureComplexGraph") { HIP_CHECK(hipStreamSynchronize(streams[0])); for (size_t i = 0; i < N; i++) { if (i > (N / 2 - 1)) { - REQUIRE(Ch.host_ptr()[i] == (Bh.host_ptr()[i] + Ah.host_ptr()[i] * CONST_KER2_VAL)); + REQUIRE(Ch.host_ptr()[i] == (Bh.host_ptr()[i] + + Ah.host_ptr()[i] * CONST_KER2_VAL)); } else { - REQUIRE(Ch.host_ptr()[i] == - (Bh.host_ptr()[i] + Ah.host_ptr()[i] * CONST_KER1_VAL * CONST_KER3_VAL)); + REQUIRE(Ch.host_ptr()[i] == (Bh.host_ptr()[i] + + Ah.host_ptr()[i] * CONST_KER1_VAL * CONST_KER3_VAL)); } } } @@ -1267,3 +1317,226 @@ TEST_CASE("Unit_hipStreamBeginCapture_Positive_captureEmptyStreams") { HIP_CHECK(hipGraphDestroy(graph)); } + +/** + * Test Description + * ------------------------ + * - Test to verify hipStreamSynchronize on a stream works when stream capture + * on another stream is ongoing. + * Test source + * ------------------------ + * - catch\unit\graph\hipStreamBeginCapture.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.6 + */ + +TEST_CASE("Unit_hipStreamBeginCapture_StreamSync_OngoingCapture") { + hipStreamCaptureMode flag = hipStreamCaptureModeRelaxed; + constexpr int GRIDSIZE = 1; + constexpr int BLOCKSIZE = 512; + constexpr int VALUE1 = 7, VALUE2 = 11; + hipGraph_t graph{nullptr}; + hipGraphExec_t graphExec{nullptr}; + // Allocate device memory + LinearAllocGuard Ah = LinearAllocGuard(LinearAllocs::malloc, + BLOCKSIZE * sizeof(int)); + LinearAllocGuard Ad = LinearAllocGuard(LinearAllocs::hipMalloc, + BLOCKSIZE * sizeof(int)); + LinearAllocGuard Bh = LinearAllocGuard(LinearAllocs::malloc, + BLOCKSIZE * sizeof(int)); + LinearAllocGuard Bd = LinearAllocGuard(LinearAllocs::hipMalloc, + BLOCKSIZE * sizeof(int)); + // Fill input data + std::fill_n(Ah.host_ptr(), BLOCKSIZE, VALUE1); + std::fill_n(Bh.host_ptr(), BLOCKSIZE, VALUE2); + // Stream create + StreamsGuard stream0(1); + // Capture streams into graph + SECTION("Stream Creation Before Capture") { + StreamsGuard stream1(1); + HIP_CHECK(hipStreamBeginCapture(stream0[0], flag)); + HIP_CHECK(hipMemcpyAsync(Ad.ptr(), Ah.host_ptr(), BLOCKSIZE * sizeof(int), + hipMemcpyDefault, stream1[0])); + HIP_CHECK(hipMemcpyAsync(Bd.ptr(), Bh.host_ptr(), BLOCKSIZE * sizeof(int), + hipMemcpyDefault, stream1[0])); + HIP_CHECK(hipStreamSynchronize(stream1[0])); + myadd<<>>(Ad.ptr(), Bd.ptr()); + HIP_CHECK(hipStreamEndCapture(stream0[0], &graph)); // End Capture + } + SECTION("Synchronizing multiple streams during Capture") { + StreamsGuard stream1(1), stream2(1); + HIP_CHECK(hipStreamBeginCapture(stream0[0], flag)); + HIP_CHECK(hipMemcpyAsync(Ad.ptr(), Ah.host_ptr(), BLOCKSIZE * sizeof(int), + hipMemcpyDefault, stream1[0])); + HIP_CHECK(hipMemcpyAsync(Bd.ptr(), Bh.host_ptr(), BLOCKSIZE * sizeof(int), + hipMemcpyDefault, stream2[0])); + HIP_CHECK(hipStreamSynchronize(stream1[0])); + HIP_CHECK(hipStreamSynchronize(stream2[0])); + myadd<<>>(Ad.ptr(), Bd.ptr()); + HIP_CHECK(hipStreamEndCapture(stream0[0], &graph)); // End Capture + } + SECTION("Stream Creation After Capture") { + HIP_CHECK(hipStreamBeginCapture(stream0[0], flag)); + StreamsGuard stream1(1); + HIP_CHECK(hipMemcpyAsync(Ad.ptr(), Ah.host_ptr(), BLOCKSIZE * sizeof(int), + hipMemcpyDefault, stream1[0])); + HIP_CHECK(hipMemcpyAsync(Bd.ptr(), Bh.host_ptr(), BLOCKSIZE * sizeof(int), + hipMemcpyDefault, stream1[0])); + HIP_CHECK(hipStreamSynchronize(stream1[0])); + myadd<<>>(Ad.ptr(), Bd.ptr()); + HIP_CHECK(hipStreamEndCapture(stream0[0], &graph)); // End Capture + } + SECTION("Stream Synchronize Before Capture") { + StreamsGuard stream1(1); + HIP_CHECK(hipMemcpyAsync(Ad.ptr(), Ah.host_ptr(), BLOCKSIZE * sizeof(int), + hipMemcpyDefault, stream1[0])); + HIP_CHECK(hipMemcpyAsync(Bd.ptr(), Bh.host_ptr(), BLOCKSIZE * sizeof(int), + hipMemcpyDefault, stream1[0])); + HIP_CHECK(hipStreamSynchronize(stream1[0])); + HIP_CHECK(hipStreamBeginCapture(stream0[0], flag)); + myadd<<>>(Ad.ptr(), Bd.ptr()); + HIP_CHECK(hipStreamEndCapture(stream0[0], &graph)); // End Capture + } + SECTION("Stream Synchronize After Capture") { + HIP_CHECK(hipStreamBeginCapture(stream0[0], flag)); + myadd<<>>(Ad.ptr(), Bd.ptr()); + HIP_CHECK(hipStreamEndCapture(stream0[0], &graph)); // End Capture + StreamsGuard stream1(1); + HIP_CHECK(hipMemcpyAsync(Ad.ptr(), Ah.host_ptr(), BLOCKSIZE * sizeof(int), + hipMemcpyDefault, stream1[0])); + HIP_CHECK(hipMemcpyAsync(Bd.ptr(), Bh.host_ptr(), BLOCKSIZE * sizeof(int), + hipMemcpyDefault, stream1[0])); + HIP_CHECK(hipStreamSynchronize(stream1[0])); + } + // Execute and test the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, stream0[0])); + HIP_CHECK(hipStreamSynchronize(stream0[0])); + // Check output + HIP_CHECK(hipMemcpy(Ah.host_ptr(), Ad.ptr(), BLOCKSIZE * sizeof(int), + hipMemcpyDeviceToHost)); + for (int idx = 0; idx < BLOCKSIZE; idx++) { + REQUIRE(Ah.host_ptr()[idx] == (VALUE1 + VALUE2)); + } + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); +} + +/** + * Test Description + * ------------------------ + * - Test to verify hipStreamSynchronize on a stream behavior when stream capture + * on another stream is ongoing in another thread. + * Test source + * ------------------------ + * - catch\unit\graph\hipStreamBeginCapture.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.6 + */ +// Local function executed as thread +static void strmSyncThread(int *Ah, int *Ad, int *Bh, int *Bd, + int BLOCKSIZE, hipError_t *error) { + StreamsGuard stream(1); + HIP_CHECK(hipMemcpyAsync(Ad, Ah, BLOCKSIZE * sizeof(int), + hipMemcpyDefault, stream[0])); + HIP_CHECK(hipMemcpyAsync(Bd, Bh, BLOCKSIZE * sizeof(int), + hipMemcpyDefault, stream[0])); + *error = hipStreamSynchronize(stream[0]); +} + +// Local function executed as thread +static void captureStrmThread(hipGraph_t *graph, int *Ah, int *Ad, + int *Bh, int *Bd, int BLOCKSIZE, int GRIDSIZE, + hipStreamCaptureMode flag, hipError_t *error) { + StreamsGuard stream(1); + // Capture streams into graph + HIP_CHECK(hipStreamBeginCapture(stream[0], flag)); + std::thread t1(strmSyncThread, Ah, Ad, Bh, Bd, BLOCKSIZE, error); + t1.join(); + myadd<<>>(Ad, Bd); + HIP_CHECK(hipStreamEndCapture(stream[0], graph)); // End Capture +} + +TEST_CASE("Unit_hipStreamBeginCapture_StreamSync_OngoingCapture_MThread") { + constexpr int GRIDSIZE = 1; + constexpr int BLOCKSIZE = 512; + constexpr int VALUE1 = 7, VALUE2 = 11; + hipGraph_t graph{nullptr}; + // Allocate device memory + LinearAllocGuard Ah = LinearAllocGuard(LinearAllocs::malloc, + BLOCKSIZE * sizeof(int)); + LinearAllocGuard Ad = LinearAllocGuard(LinearAllocs::hipMalloc, + BLOCKSIZE * sizeof(int)); + LinearAllocGuard Bh = LinearAllocGuard(LinearAllocs::malloc, + BLOCKSIZE * sizeof(int)); + LinearAllocGuard Bd = LinearAllocGuard(LinearAllocs::hipMalloc, + BLOCKSIZE * sizeof(int)); + // Fill input data + std::fill_n(Ah.host_ptr(), BLOCKSIZE, VALUE1); + std::fill_n(Bh.host_ptr(), BLOCKSIZE, VALUE2); + // Stream create + hipError_t error = hipSuccess; + SECTION("Capture Flag = hipStreamCaptureModeGlobal Single Threaded") { + StreamsGuard stream(2); + // Capture streams into graph + HIP_CHECK(hipStreamBeginCapture(stream[0], hipStreamCaptureModeGlobal)); + HIP_CHECK(hipMemcpyAsync(Ad.ptr(), Ah.host_ptr(), + BLOCKSIZE * sizeof(int), hipMemcpyDefault, stream[1])); + HIP_CHECK(hipMemcpyAsync(Bd.ptr(), Bh.host_ptr(), + BLOCKSIZE * sizeof(int), hipMemcpyDefault, stream[1])); + error = hipStreamSynchronize(stream[1]); + REQUIRE(error == hipErrorStreamCaptureUnsupported); + } +#if HT_NVIDIA + SECTION("Capture Flag = hipStreamCaptureModeThreadLocal Single Threaded") { + StreamsGuard stream(2); + // Capture streams into graph + HIP_CHECK(hipStreamBeginCapture(stream[0], + hipStreamCaptureModeThreadLocal)); + HIP_CHECK(hipMemcpyAsync(Ad.ptr(), Ah.host_ptr(), + BLOCKSIZE * sizeof(int), hipMemcpyDefault, stream[1])); + HIP_CHECK(hipMemcpyAsync(Bd.ptr(), Bh.host_ptr(), + BLOCKSIZE * sizeof(int), hipMemcpyDefault, stream[1])); + error = hipStreamSynchronize(stream[1]); + REQUIRE(error == hipErrorStreamCaptureUnsupported); + } +#endif +#if HT_AMD + SECTION("Capture Flag = hipStreamCaptureModeGlobal Multithreaded") { + captureStrmThread(&graph, Ah.host_ptr(), Ad.ptr(), + Bh.host_ptr(), Bd.ptr(), BLOCKSIZE, GRIDSIZE, + hipStreamCaptureModeGlobal, &error); + REQUIRE(error == hipErrorStreamCaptureUnsupported); + } +#endif + SECTION("Capture Flag = hipStreamCaptureModeThreadLocal Multithreaded") { + captureStrmThread(&graph, Ah.host_ptr(), Ad.ptr(), + Bh.host_ptr(), Bd.ptr(), BLOCKSIZE, GRIDSIZE, + hipStreamCaptureModeThreadLocal, &error); + REQUIRE(error == hipSuccess); + } + SECTION("Capture Flag = hipStreamCaptureModeRelaxed Multithreaded") { + captureStrmThread(&graph, Ah.host_ptr(), Ad.ptr(), + Bh.host_ptr(), Bd.ptr(), BLOCKSIZE, GRIDSIZE, + hipStreamCaptureModeRelaxed, &error); + REQUIRE(error == hipSuccess); + } + if (graph != nullptr) { + hipGraphExec_t graphExec{nullptr}; + StreamsGuard stream(1); + // Execute and test the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, stream[0])); + HIP_CHECK(hipStreamSynchronize(stream[0])); + // Check output + HIP_CHECK(hipMemcpy(Ah.host_ptr(), Ad.ptr(), BLOCKSIZE * sizeof(int), + hipMemcpyDeviceToHost)); + for (int idx = 0; idx < BLOCKSIZE; idx++) { + REQUIRE(Ah.host_ptr()[idx] == (VALUE1 + VALUE2)); + } + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + } +}