From 17d3f51effbaaf668fa2b790e68c273fd5257e13 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Wed, 21 Jun 2023 11:53:18 +0530 Subject: [PATCH] SWDEV-372504 - [catch2][dtest] Added additional functional tests for the API hipGraphAddMemcpyNode() and hipGraphAddMemcpyNode1D() (#146) Change-Id: I5bd0f8540adf4938606d650863e5e31e6a6aaff8 --- catch/unit/graph/hipGraphAddMemcpyNode.cc | 52 ++++++++++++++- catch/unit/graph/hipGraphAddMemcpyNode1D.cc | 72 ++++++++++----------- 2 files changed, 87 insertions(+), 37 deletions(-) diff --git a/catch/unit/graph/hipGraphAddMemcpyNode.cc b/catch/unit/graph/hipGraphAddMemcpyNode.cc index 5612f998ef..745b29d41a 100644 --- a/catch/unit/graph/hipGraphAddMemcpyNode.cc +++ b/catch/unit/graph/hipGraphAddMemcpyNode.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 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 @@ -48,10 +48,15 @@ Testcase Scenarios : Functional Memcpy nodes are added and assigned to Peer device. 4) Perform memcpy operation for 1D, 2D and 3D arrays on Peer device and verify the results. +5) Create two host pointers, copy the data between them by the api + hipGraphAddMemcpyNode with data transfer kind hipMemcpyHostToHost. + Validate the output. */ #include #include +#include +#include #define ZSIZE 32 #define YSIZE 32 @@ -517,4 +522,49 @@ TEST_CASE("Unit_hipGraphAddMemcpyNode_PeerAccessFunctional") { validateMemcpyNode1DArray(true); } } +/* +* Create two host pointers, copy the data between them by the api +* hipGraphAddMemcpyNode with data transfer kind hipMemcpyHostToHost. +* Validate the output. +*/ +TEST_CASE("Unit_hipGraphAddMemcpyNode_HostToHost") { + constexpr size_t size = 1024; + size_t numW = size * sizeof(int); + // Host Vectors + std::vector A_h(numW); + std::vector B_h(numW); + // Initialization + std::iota(A_h.begin(), A_h.end(), 0); + std::fill_n(B_h.begin(), size, 0); + hipGraph_t graph; + hipStream_t streamForGraph; + hipGraphExec_t graphExec; + hipGraphNode_t memcpyH2H; + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipStreamCreate(&streamForGraph)); + + hipMemcpy3DParms myparms{}; + myparms.srcPos = make_hipPos(0, 0, 0); + myparms.dstPos = make_hipPos(0, 0, 0); + myparms.srcPtr = make_hipPitchedPtr(A_h.data(), numW, numW, 1); + myparms.dstPtr = make_hipPitchedPtr(B_h.data(), numW, numW, 1); + myparms.extent = make_hipExtent(numW, 1, 1); + myparms.kind = hipMemcpyHostToHost; + + // Host to Host + HIP_CHECK(hipGraphAddMemcpyNode(&memcpyH2H, graph, nullptr, + 0, &myparms)); + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + + // Validation + REQUIRE(memcmp(A_h.data(), B_h.data(), numW) == 0); +} diff --git a/catch/unit/graph/hipGraphAddMemcpyNode1D.cc b/catch/unit/graph/hipGraphAddMemcpyNode1D.cc index e68eac22cf..9ed4a9fbcd 100644 --- a/catch/unit/graph/hipGraphAddMemcpyNode1D.cc +++ b/catch/unit/graph/hipGraphAddMemcpyNode1D.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 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 @@ -24,6 +24,9 @@ Functional - Memcpy nodes are added and assigned to default device. 2) Allocate memory on default device(Dev 0), Perform memcpy operation for 1D arrays on Peer device(Dev 1) and verify the results. +3) Create two host pointers, copy the data between them by the api hipGraphAddMemcpyNode1D with data transfer + kind hipMemcpyHostToHost. Validate the output. + Negative - 1) Pass pGraphNode as nullptr and check if api returns error. 2) When graph is un-initialized argument(skipping graph creation), api should return error code. @@ -38,7 +41,8 @@ Negative - #include #include - +#include +#include static void validateMemcpyNode1DArray(bool peerAccess) { constexpr int SIZE{32}; @@ -197,46 +201,42 @@ TEST_CASE("Unit_hipGraphAddMemcpyNode1D_Negative") { HIP_CHECK(hipFree(A_h)); HIP_CHECK(hipGraphDestroy(graph)); } +/* + * Create two host pointers, copy the data between them by the api + * hipGraphAddMemcpyNode1D with data transfer kind hipMemcpyHostToHost. + * Validate the output. +*/ +TEST_CASE("Unit_hipGraphAddMemcpyNode1D_HostToHost") { + constexpr size_t size = 1024; + size_t numBytes{size * sizeof(int)}; -TEST_CASE("Unit_hipGraphAddMemcpyNode1D_Negative_Basic") { - constexpr size_t N = 1024; - constexpr size_t Nbytes = N * sizeof(int); - int *A_d, *B_d, *C_d, *A_h, *B_h, *C_h; + // Host Vectors + std::vector A_h(size); + std::vector B_h(size); + // Initialization + std::iota(A_h.begin(), A_h.end(), 0); + std::fill_n(B_h.begin(), size, 0); - hipGraphNode_t memcpy_A, memcpy_B, memcpy_C; - hipError_t ret; hipGraph_t graph; - - HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); - + hipStream_t streamForGraph; + hipGraphExec_t graphExec; + hipGraphNode_t memcpyH2H; HIP_CHECK(hipGraphCreate(&graph, 0)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpy_A, graph, nullptr, 0, A_d, A_h, - Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipStreamCreate(&streamForGraph)); - // Pass memcpy direction as hipMemcpyDeviceToHost and - // source pointer as host and destination pointer as device pointer - ret = hipGraphAddMemcpyNode1D(&memcpy_B, graph, nullptr, 0, B_d, B_h, - Nbytes, hipMemcpyDeviceToHost); - REQUIRE(hipErrorInvalidValue == ret); + // Host to Host + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2H, graph, nullptr, 0, + B_h.data(), A_h.data(), numBytes, hipMemcpyHostToHost)); - // Pass memcpy direction as hipMemcpyHostToDevice and - // source pointer as device and destination pointer as host pointer - ret = hipGraphAddMemcpyNode1D(&memcpy_C, graph, nullptr, 0, C_h, C_d, - Nbytes, hipMemcpyHostToDevice); - REQUIRE(hipErrorInvalidValue == ret); + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); - // Pass memcpy direction as hipMemcpyDeviceToDevice and - // pass source pointer as device and destination pointer as host pointer - ret = hipGraphAddMemcpyNode1D(&memcpy_C, graph, nullptr, 0, C_h, C_d, - Nbytes, hipMemcpyDeviceToDevice); - REQUIRE(hipErrorInvalidValue == ret); - - // Pass memcpy direction as hipMemcpyDeviceToDevice and - // pass source pointer as host and destination pointer as device pointer - ret = hipGraphAddMemcpyNode1D(&memcpy_C, graph, nullptr, 0, C_d, C_h, - Nbytes, hipMemcpyDeviceToDevice); - REQUIRE(hipErrorInvalidValue == ret); - - HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + + // Validation + REQUIRE(std::equal(A_h.begin(), A_h.end(), B_h.begin(), B_h.end())); }