diff --git a/projects/hip/tests/catch/unit/graph/CMakeLists.txt b/projects/hip/tests/catch/unit/graph/CMakeLists.txt index 439ac1ffed..7f15b7e6ff 100644 --- a/projects/hip/tests/catch/unit/graph/CMakeLists.txt +++ b/projects/hip/tests/catch/unit/graph/CMakeLists.txt @@ -79,6 +79,7 @@ set(TEST_SRC hipGraphNodeGetDependencies.cc hipGraphHostNodeGetParams.cc hipGraphExecChildGraphNodeSetParams.cc + hipStreamGetCaptureInfo_v2.cc ) hip_add_exe_to_target(NAME GraphsTest diff --git a/projects/hip/tests/catch/unit/graph/hipStreamGetCaptureInfo_v2.cc b/projects/hip/tests/catch/unit/graph/hipStreamGetCaptureInfo_v2.cc new file mode 100644 index 0000000000..d19e059083 --- /dev/null +++ b/projects/hip/tests/catch/unit/graph/hipStreamGetCaptureInfo_v2.cc @@ -0,0 +1,280 @@ +/* +Copyright (c) 2022 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 +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/** + +Testcase Scenarios +------------------ +Functional: +1) Start stream capture and get capture info v2. Verify api is success, capture status is hipStreamCaptureStatusActive, + identifier returned is valid/non-zero, graph object is returned. +2) When stream capture is in progress, create dependent nodes by creating multistream dependencies and verify the api returns + valid dependent nodes. +3) End stream capture and get capture info. Verify api is success, capture status is hipStreamCaptureStatusNone and + identifier/graph/nodes are not returned by api. +4) When optional parameters are not passed, make sure api still returns capture status of stream. +5) Begin capture on hipStreamPerThread, get capture info v2 and validate results. +6) Perform multiple captures and verify the identifier returned is unique. + +Parameter Validation/Negative: +1) Capture Status location as nullptr and verify api returns error code. +2) Stream as nullptr and verify api returns error code. + +*/ + +#include +#include +#include + +constexpr size_t N = 1000000; +constexpr int LAUNCH_ITERS = 1; + +/** + * Validates stream capture infov2, launches graph and verifies results + */ +void validateStreamCaptureInfoV2(hipStream_t mstream) { + hipStream_t stream1{nullptr}, stream2{nullptr}, streamForLaunch{nullptr}; + hipEvent_t memcpyEvent1, memsetEvent2, forkStreamEvent; + hipGraph_t graph{nullptr}, capInfoGraph{nullptr}; + hipGraphExec_t graphExec{nullptr}; + constexpr unsigned blocks = 512; + constexpr unsigned threadsPerBlock = 256; + const hipGraphNode_t *nodelist{}; + size_t Nbytes = N * sizeof(float), numDependencies; + float *A_d, *C_d; + float *A_h, *C_h; + A_h = reinterpret_cast(malloc(Nbytes)); + C_h = reinterpret_cast(malloc(Nbytes)); + REQUIRE(A_h != nullptr); + REQUIRE(C_h != nullptr); + HIP_CHECK(hipMalloc(&A_d, Nbytes)); + HIP_CHECK(hipMalloc(&C_d, Nbytes)); + REQUIRE(A_d != nullptr); + REQUIRE(C_d != nullptr); + HIP_CHECK(hipStreamCreate(&streamForLaunch)); + + // Initialize input buffer + for (size_t i = 0; i < N; ++i) { + A_h[i] = 3.146f + i; // Pi + } + + // Create cross stream dependencies. + // memset/memcpy operations are done on stream1 and stream2 + // and they are joined back to mainstream + HIP_CHECK(hipStreamCreate(&stream1)); + HIP_CHECK(hipStreamCreate(&stream2)); + HIP_CHECK(hipEventCreate(&memcpyEvent1)); + HIP_CHECK(hipEventCreate(&memsetEvent2)); + HIP_CHECK(hipEventCreate(&forkStreamEvent)); + + HIP_CHECK(hipStreamBeginCapture(mstream, hipStreamCaptureModeGlobal)); + HIP_CHECK(hipEventRecord(forkStreamEvent, mstream)); + HIP_CHECK(hipStreamWaitEvent(stream1, forkStreamEvent, 0)); + HIP_CHECK(hipStreamWaitEvent(stream2, forkStreamEvent, 0)); + HIP_CHECK(hipMemsetAsync(A_d, 0, Nbytes, stream1)); + HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, stream1)); + HIP_CHECK(hipEventRecord(memcpyEvent1, stream1)); + HIP_CHECK(hipMemsetAsync(C_d, 0, Nbytes, stream2)); + HIP_CHECK(hipEventRecord(memsetEvent2, stream2)); + HIP_CHECK(hipStreamWaitEvent(mstream, memcpyEvent1, 0)); + HIP_CHECK(hipStreamWaitEvent(mstream, memsetEvent2, 0)); + + hipStreamCaptureStatus captureStatus{hipStreamCaptureStatusNone}; + unsigned long long capSequenceID = 0; // NOLINT + constexpr int numDepsCreated = 2; // Num of dependencies created + + HIP_CHECK(hipStreamGetCaptureInfo_v2(mstream, &captureStatus, + &capSequenceID, &capInfoGraph, &nodelist, &numDependencies)); + + // verify capture status is active, sequence id is valid, graph is returned, + REQUIRE(captureStatus == hipStreamCaptureStatusActive); + REQUIRE(capSequenceID > 0); + REQUIRE(capInfoGraph != nullptr); + REQUIRE(numDependencies == numDepsCreated); + + // verify dependency nodes list returned is the one we created. + hipGraphNodeType type(hipGraphNodeTypeEmpty); + + HIP_CHECK(hipGraphNodeGetType(nodelist[0], &type)); + if ((type != hipGraphNodeTypeMemset) && (type != hipGraphNodeTypeMemcpy)) { + INFO("Type0 returned as " << type); + REQUIRE(false); + } + + HIP_CHECK(hipGraphNodeGetType(nodelist[1], &type)); + if ((type != hipGraphNodeTypeMemset) && (type != hipGraphNodeTypeMemcpy)) { + INFO("Type1 returned as " << type); + REQUIRE(false); + } + + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), + dim3(threadsPerBlock), 0, mstream, A_d, C_d, N); + + // End capture and verify graph is returned + HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, mstream)); + HIP_CHECK(hipStreamEndCapture(mstream, &graph)); + REQUIRE(graph != nullptr); + + // verify capture status is inactive and other params are not updated + capSequenceID = 0; + capInfoGraph = nullptr; + numDependencies = 0; + nodelist = nullptr; + HIP_CHECK(hipStreamGetCaptureInfo_v2(mstream, &captureStatus, + &capSequenceID, &capInfoGraph, &nodelist, &numDependencies)); + REQUIRE(captureStatus == hipStreamCaptureStatusNone); + REQUIRE(capSequenceID == 0); + REQUIRE(capInfoGraph == nullptr); + REQUIRE(nodelist == nullptr); + REQUIRE(numDependencies == 0); + + // Verify api still returns capture status when optional args are not passed + HIP_CHECK(hipStreamGetCaptureInfo_v2(mstream, &captureStatus)); + REQUIRE(captureStatus == hipStreamCaptureStatusNone); + + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + REQUIRE(graphExec != nullptr); + + // Replay the recorded sequence multiple times + for (int i = 0; i < LAUNCH_ITERS; i++) { + HIP_CHECK(hipGraphLaunch(graphExec, streamForLaunch)); + } + + HIP_CHECK(hipStreamSynchronize(streamForLaunch)); + + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForLaunch)); + HIP_CHECK(hipStreamDestroy(stream1)); + HIP_CHECK(hipStreamDestroy(stream2)); + HIP_CHECK(hipEventDestroy(forkStreamEvent)); + HIP_CHECK(hipEventDestroy(memcpyEvent1)); + HIP_CHECK(hipEventDestroy(memsetEvent2)); + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipFree(C_d)); + + // Validate the computation + for (size_t i = 0; i < N; i++) { + if (C_h[i] != A_h[i] * A_h[i]) { + INFO("A and C not matching at " << i << " C_h[i] " << C_h[i] + << " A_h[i] " << A_h[i]); + REQUIRE(false); + } + } + free(A_h); + free(C_h); +} + +/** + * Basic Functional Test for stream capture and getting capture info V2. + * Regular/custom stream is used for stream capture. + */ +TEST_CASE("Unit_hipStreamGetCaptureInfo_v2_BasicFunctional") { + hipStream_t streamForCapture; + + HIP_CHECK(hipStreamCreate(&streamForCapture)); + validateStreamCaptureInfoV2(streamForCapture); + HIP_CHECK(hipStreamDestroy(streamForCapture)); +} + +/** + * Test performs stream capture on hipStreamPerThread and validates + * capture info V2. + */ +TEST_CASE("Unit_hipStreamGetCaptureInfo_v2_hipStreamPerThread") { + validateStreamCaptureInfoV2(hipStreamPerThread); +} + +/** + * Test starts stream capture on multiple streams and verifies uniqueness of + * identifiers returned from capture Info V2. + */ +TEST_CASE("Unit_hipStreamGetCaptureInfo_v2_UniqueID") { + constexpr int numStreams = 100; + hipStream_t streams[numStreams]{}; + hipStreamCaptureStatus captureStatus{hipStreamCaptureStatusNone}; + std::vector idlist; + unsigned long long capSequenceID{}; // NOLINT + hipGraph_t graph{nullptr}; + + for (int i = 0; i < numStreams; i++) { + HIP_CHECK(hipStreamCreate(&streams[i])); + HIP_CHECK(hipStreamBeginCapture(streams[i], hipStreamCaptureModeGlobal)); + HIP_CHECK(hipStreamGetCaptureInfo_v2(streams[i], &captureStatus, + &capSequenceID, nullptr, nullptr, nullptr)); + REQUIRE(captureStatus == hipStreamCaptureStatusActive); + REQUIRE(capSequenceID > 0); + idlist.push_back(capSequenceID); + } + + for (int i = 0; i < numStreams; i++) { + for (int j = i+1; j < numStreams; j++) { + if (idlist[i] == idlist[j]) { + INFO("Same identifier returned for stream " + << i << " and stream " << j); + REQUIRE(false); + } + } + } + + for (int i = 0; i < numStreams; i++) { + HIP_CHECK(hipStreamEndCapture(streams[i], &graph)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streams[i])); + } +} + +/** + * Parameter validation/Negative tests for api + */ +TEST_CASE("Unit_hipStreamGetCaptureInfo_v2_ParamValidation") { + hipError_t ret; + hipStream_t stream; + float *A_d; + hipGraph_t graph{}, capInfoGraph{}; + hipStreamCaptureStatus captureStatus; + unsigned long long capSequenceID; // NOLINT + size_t numDependencies; + const hipGraphNode_t *nodelist{}; + constexpr int numBytes{100}; + + HIP_CHECK(hipMalloc(&A_d, numBytes)); + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal)); + HIP_CHECK(hipMemsetAsync(A_d, 0, numBytes, stream)); + + SECTION("Capture Status location as nullptr") { + ret = hipStreamGetCaptureInfo_v2(stream, nullptr, + &capSequenceID, &capInfoGraph, &nodelist, &numDependencies); + REQUIRE(ret == hipErrorInvalidValue); + } + + SECTION("Stream as nullptr") { + ret = hipStreamGetCaptureInfo_v2(nullptr, &captureStatus, + &capSequenceID, &capInfoGraph, &nodelist, &numDependencies); + if ((ret != hipErrorUnknown) && (ret != hipErrorStreamCaptureImplicit)) { + INFO("Ret : " << ret); + REQUIRE(false); + } + } + + HIP_CHECK(hipStreamEndCapture(stream, &graph)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipFree(A_d)); +}