From 2fcb8178c836b87614aac169b49c67017ed381b2 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Wed, 2 Nov 2022 14:42:11 +0530 Subject: [PATCH] SWDEV-306122 - [catch2][dtest] Additional tests for hipGraphGetNodes/hipGraphGetRootNodes api (#2860) Change-Id: I4e9a6eadd92b6de4aded0aa96464bfcc59441ac4 [ROCm/hip-tests commit: 664c8b377d5443554a3df93f0772ce7175963a92] --- .../catch/unit/graph/hipGraphGetNodes.cc | 106 +++++++++++++++++ .../catch/unit/graph/hipGraphGetRootNodes.cc | 109 ++++++++++++++++++ 2 files changed, 215 insertions(+) diff --git a/projects/hip-tests/catch/unit/graph/hipGraphGetNodes.cc b/projects/hip-tests/catch/unit/graph/hipGraphGetNodes.cc index 18b3df1722..93d86ac87f 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphGetNodes.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphGetNodes.cc @@ -25,6 +25,8 @@ Functional :: 2) Pass nodes as nullptr and verify numNodes returns actual number of nodes added to graph. 3) If numNodes passed is greater than the actual number of nodes, the remaining entries in nodes will be set to NULL, and the number of nodes actually obtained will be returned in numNodes. +4) Begin stream capture and push operations to stream. Verify nodes of created graph are matching the +operations pushed. Argument Validation :: 1) Pass graph as nullptr and verify api returns error code. @@ -139,6 +141,110 @@ TEST_CASE("Unit_hipGraphGetNodes_Functional") { free(nodes); } +/** + * Begin stream capture and push operations to stream. + * Verify nodes of created graph are matching the operations pushed. + */ +TEST_CASE("Unit_hipGraphGetNodes_CapturedStream") { + hipGraph_t graph{nullptr}; + hipGraphExec_t graphExec{nullptr}; + constexpr unsigned blocks = 512; + constexpr unsigned threadsPerBlock = 256; + constexpr size_t N = 1000000; + size_t Nbytes = N * sizeof(float); + constexpr int numMemcpy{2}, numKernel{1}, numMemset{1}; + int cntMemcpy{}, cntKernel{}, cntMemset{}; + hipStream_t stream, streamForGraph; + hipGraphNodeType nodeType; + 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(&streamForGraph)); + // Initialize input buffer + for (size_t i = 0; i < N; ++i) { + A_h[i] = 3.146f + i; // Pi + } + + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipStreamBeginCapture(stream, hipStreamCaptureModeGlobal)); + 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); + HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamEndCapture(stream, &graph)); + REQUIRE(graph != nullptr); + + size_t numNodes{}; + HIP_CHECK(hipGraphGetNodes(graph, nullptr, &numNodes)); + INFO("Num of nodes returned by GetNodes : " << numNodes); + REQUIRE(numNodes == numMemcpy + numKernel + numMemset); + + int numBytes = sizeof(hipGraphNode_t) * numNodes; + hipGraphNode_t* nodes = reinterpret_cast(malloc(numBytes)); + REQUIRE(nodes != nullptr); + + HIP_CHECK(hipGraphGetNodes(graph, nodes, &numNodes)); + for (size_t i = 0; i < numNodes; i++) { + HIP_CHECK(hipGraphNodeGetType(nodes[i], &nodeType)); + + switch (nodeType) { + case hipGraphNodeTypeMemcpy: + cntMemcpy++; + break; + + case hipGraphNodeTypeKernel: + cntKernel++; + break; + + case hipGraphNodeTypeMemset: + cntMemset++; + break; + + default: + INFO("Unexpected nodetype returned : " << nodeType); + REQUIRE(false); + } + } + + REQUIRE(cntMemcpy == numMemcpy); + REQUIRE(cntKernel == numKernel); + REQUIRE(cntMemset == numMemset); + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // 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); + } + } + + HIP_CHECK(hipStreamDestroy(streamForGraph)); + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + free(A_h); + free(C_h); + free(nodes); + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipFree(C_d)); +} + /** * Test performs api parameter validation by passing various values * as input and output parameters and validates the behavior. diff --git a/projects/hip-tests/catch/unit/graph/hipGraphGetRootNodes.cc b/projects/hip-tests/catch/unit/graph/hipGraphGetRootNodes.cc index c2484fc8b5..2be8ea5110 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphGetRootNodes.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphGetRootNodes.cc @@ -26,6 +26,8 @@ Functional :: 2) Pass nodes as nullptr and verify api returns actual number of root nodes added to graph. 3) If NumRootNodes passed is greater than the actual number of root nodes, the remaining entries in nodes list will be set to NULL, and the number of nodes actually obtained will be returned in NumRootNodes. + 4) Create a graph with stream capture done on multiple dependent streams. + Verify root nodes of created graph are matching the operations pushed which doesn't have dependencies. Argument Validation :: 1) Pass graph as nullptr and verify api returns error code. @@ -50,6 +52,7 @@ TEST_CASE("Unit_hipGraphGetRootNodes_Functional") { constexpr auto addlEntries = 5; hipGraph_t graph; + hipGraphNode_t memcpyNode, kernelNode; hipKernelNodeParams kernelNodeParams{}; hipStream_t streamForGraph; @@ -141,6 +144,112 @@ TEST_CASE("Unit_hipGraphGetRootNodes_Functional") { free(rootnodes); } + +/** + * Create a graph with stream capture done on multiple dependent streams. Verify root nodes + * of created graph are matching the operations pushed which doesn't have dependencies. + */ +TEST_CASE("Unit_hipGraphGetRootNodes_CapturedStream") { + hipStream_t stream1{nullptr}, stream2{nullptr}, mstream{nullptr}; + hipStream_t streamForGraph{nullptr}; + hipEvent_t memsetEvent1, memsetEvent2, forkStreamEvent; + hipGraph_t graph{nullptr}; + hipGraphExec_t graphExec{nullptr}; + constexpr unsigned blocks = 512; + constexpr unsigned threadsPerBlock = 256; + constexpr size_t N = 1000000; + constexpr int numMemsetNodes = 2; + size_t Nbytes = N * sizeof(float), numRootNodes{}; + 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(&streamForGraph)); + + // Initialize input buffer + for (size_t i = 0; i < N; ++i) { + A_h[i] = 3.146f + i; // Pi + } + + HIP_CHECK(hipStreamCreate(&stream1)); + HIP_CHECK(hipStreamCreate(&stream2)); + HIP_CHECK(hipStreamCreate(&mstream)); + HIP_CHECK(hipEventCreate(&memsetEvent1)); + 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(hipEventRecord(memsetEvent1, stream1)); + HIP_CHECK(hipMemsetAsync(C_d, 0, Nbytes, stream2)); + HIP_CHECK(hipEventRecord(memsetEvent2, stream2)); + HIP_CHECK(hipStreamWaitEvent(mstream, memsetEvent1, 0)); + HIP_CHECK(hipStreamWaitEvent(mstream, memsetEvent2, 0)); + HIP_CHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, mstream)); + hipLaunchKernelGGL(HipTest::vector_square, dim3(blocks), + dim3(threadsPerBlock), 0, mstream, A_d, C_d, N); + HIP_CHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, mstream)); + HIP_CHECK(hipStreamEndCapture(mstream, &graph)); + + // Verify numof root nodes + HIP_CHECK(hipGraphGetRootNodes(graph, nullptr, &numRootNodes)); + REQUIRE(numRootNodes == numMemsetNodes); + INFO("Num of nodes returned by GetRootNodes : " << numRootNodes); + + int numBytes = sizeof(hipGraphNode_t) * numRootNodes; + hipGraphNode_t* nodes = reinterpret_cast(malloc(numBytes)); + REQUIRE(nodes != nullptr); + + hipGraphNodeType nodeType; + HIP_CHECK(hipGraphGetRootNodes(graph, nodes, &numRootNodes)); + REQUIRE(numRootNodes == numMemsetNodes); + + // Verify root nodes returned are memset nodes. + HIP_CHECK(hipGraphNodeGetType(nodes[0], &nodeType)); + REQUIRE(nodeType == hipGraphNodeTypeMemset); + HIP_CHECK(hipGraphNodeGetType(nodes[1], &nodeType)); + REQUIRE(nodeType == hipGraphNodeTypeMemset); + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // 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); + } + } + + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + HIP_CHECK(hipStreamDestroy(mstream)); + HIP_CHECK(hipStreamDestroy(stream1)); + HIP_CHECK(hipStreamDestroy(stream2)); + HIP_CHECK(hipEventDestroy(forkStreamEvent)); + HIP_CHECK(hipEventDestroy(memsetEvent1)); + HIP_CHECK(hipEventDestroy(memsetEvent2)); + free(A_h); + free(C_h); + free(nodes); + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipFree(C_d)); +} + + /** * Test performs api parameter validation by passing various values * as input and output parameters and validates the behavior.