From 7d8d41e45bfdd49aae3d2b9a2cec09e993c7ec83 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 24 Feb 2022 18:57:27 +0530 Subject: [PATCH] SWDEV-306122 - [catch2][dtest] hipGraph tests for hipGraphAddMemcpyNode api (#2498) Change-Id: I65088f4107158e320d54b5e54a47d4ad8ce5d85f --- .../catch/unit/graph/hipGraphAddMemcpyNode.cc | 270 +++++++++++++++++- 1 file changed, 262 insertions(+), 8 deletions(-) diff --git a/tests/catch/unit/graph/hipGraphAddMemcpyNode.cc b/tests/catch/unit/graph/hipGraphAddMemcpyNode.cc index 5fcf1ccc40..6fcb3fb90e 100644 --- a/tests/catch/unit/graph/hipGraphAddMemcpyNode.cc +++ b/tests/catch/unit/graph/hipGraphAddMemcpyNode.cc @@ -19,18 +19,22 @@ THE SOFTWARE. /** Testcase Scenarios : - 1) Add multiple Memcpy nodes to graph and verify node execution is - working as expected. +1) Add memcpy node to graph and verify memcpy operation is success for all memcpy kinds(H2D, D2H and D2D). +Memcpy nodes are added and assigned to default device. +2) Perform memcpy operation for 1D, 2D and 3D arrays on default device and verify the results. +3) Add memcpy node to graph and verify memcpy operation is success for all memcpy kinds(H2D, D2H and D2D). +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. */ #include #include -/** - * Functional Test adds memcpy nodes of types H2D, D2D and D2H to graph - * and verifies execution sequence by launching graph. - */ -TEST_CASE("Unit_hipGraphAddMemcpyNode_Functional") { +#define ZSIZE 32 +#define YSIZE 32 +#define XSIZE 32 + +void validateMemcpyNode3DArray(bool peerAccess = false) { constexpr int width{10}, height{10}, depth{10}; hipArray *devArray1, *devArray2; hipChannelFormatKind formatKind = hipChannelFormatKindSigned; @@ -42,6 +46,7 @@ TEST_CASE("Unit_hipGraphAddMemcpyNode_Functional") { hipStream_t streamForGraph; hipGraphExec_t graphExec; + HIP_CHECK(hipSetDevice(0)); int *hData = reinterpret_cast(malloc(size)); int *hOutputData = reinterpret_cast(malloc(size)); @@ -69,6 +74,12 @@ TEST_CASE("Unit_hipGraphAddMemcpyNode_Functional") { make_hipExtent(width, height, depth), hipArrayDefault)); HIP_CHECK(hipGraphCreate(&graph, 0)); + // For peer access test, Memory is allocated on device(0) + // while memcpy nodes are allocated and assigned to peer device(1) + if (peerAccess) { + HIP_CHECK(hipSetDevice(1)); + } + // Host to Device memset(&myparams, 0x0, sizeof(hipMemcpy3DParms)); myparams.srcPos = make_hipPos(0, 0, 0); @@ -79,7 +90,6 @@ TEST_CASE("Unit_hipGraphAddMemcpyNode_Functional") { myparams.dstArray = devArray1; myparams.kind = hipMemcpyHostToDevice; - HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, nullptr, 0, &myparams)); dependencies.push_back(memcpyNode); @@ -126,3 +136,247 @@ TEST_CASE("Unit_hipGraphAddMemcpyNode_Functional") { free(hData); free(hOutputData); } + +void validateMemcpyNode2DArray(bool peerAccess = false) { + int harray2D[YSIZE][XSIZE]{}; + int harray2Dres[YSIZE][XSIZE]{}; + constexpr int width{XSIZE}, height{YSIZE}; + hipArray *devArray1, *devArray2; + hipChannelFormatKind formatKind = hipChannelFormatKindSigned; + hipMemcpy3DParms myparams; + hipGraph_t graph; + hipGraphNode_t memcpyNode; + std::vector dependencies; + hipStream_t streamForGraph; + hipGraphExec_t graphExec; + + HIP_CHECK(hipSetDevice(0)); + HIP_CHECK(hipStreamCreate(&streamForGraph)); + // Initialize 2D object + for (int i = 0; i < YSIZE; i++) { + for (int j = 0; j < XSIZE; j++) { + harray2D[i][j] = i + j + 1; + } + } + + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(int)*8, + 0, 0, 0, formatKind); + // Allocate 2D device array by passing depth(0) + HIP_CHECK(hipMalloc3DArray(&devArray1, &channelDesc, + make_hipExtent(width, height, 0), hipArrayDefault)); + HIP_CHECK(hipMalloc3DArray(&devArray2, &channelDesc, + make_hipExtent(width, height, 0), hipArrayDefault)); + HIP_CHECK(hipGraphCreate(&graph, 0)); + + // For peer access test, Memory is allocated on device(0) + // while memcpy nodes are allocated and assigned to peer device(1) + if (peerAccess) { + HIP_CHECK(hipSetDevice(1)); + } + + // Host to Device + memset(&myparams, 0x0, sizeof(hipMemcpy3DParms)); + myparams.srcPos = make_hipPos(0, 0, 0); + myparams.dstPos = make_hipPos(0, 0, 0); + myparams.extent = make_hipExtent(width, height, 1); + myparams.srcPtr = make_hipPitchedPtr(harray2D, width * sizeof(int), + width, height); + myparams.dstArray = devArray1; + myparams.kind = hipMemcpyHostToDevice; + + HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, nullptr, 0, &myparams)); + dependencies.push_back(memcpyNode); + + // Device to Device + memset(&myparams, 0x0, sizeof(hipMemcpy3DParms)); + myparams.srcPos = make_hipPos(0, 0, 0); + myparams.dstPos = make_hipPos(0, 0, 0); + myparams.srcArray = devArray1; + myparams.dstArray = devArray2; + myparams.extent = make_hipExtent(width, height, 1); + myparams.kind = hipMemcpyDeviceToDevice; + + HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, dependencies.data(), + dependencies.size(), &myparams)); + dependencies.clear(); + dependencies.push_back(memcpyNode); + + // Device to host + memset(&myparams, 0x0, sizeof(hipMemcpy3DParms)); + myparams.srcPos = make_hipPos(0, 0, 0); + myparams.dstPos = make_hipPos(0, 0, 0); + myparams.extent = make_hipExtent(width, height, 1); + myparams.dstPtr = make_hipPitchedPtr(harray2Dres, width * sizeof(int), + width, height); + myparams.srcArray = devArray2; + myparams.kind = hipMemcpyDeviceToHost; + + HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, dependencies.data(), + dependencies.size(), &myparams)); + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Validate result + for (int i = 0; i < YSIZE; i++) { + for (int j = 0; j < XSIZE; j++) { + if (harray2D[i][j] != harray2Dres[i][j]) { + INFO("harray2D: " << harray2D[i][j] << "harray2Dres: " + << harray2Dres[i][j] << " mismatch at (i,j) : " << i << j); + REQUIRE(false); + } + } + } + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + hipFreeArray(devArray1); + hipFreeArray(devArray2); +} + +void validateMemcpyNode1DArray(bool peerAccess = false) { + int harray1D[XSIZE]{}; + int harray1Dres[XSIZE]{}; + constexpr int width{XSIZE}; + hipArray *devArray1, *devArray2; + hipChannelFormatKind formatKind = hipChannelFormatKindSigned; + hipMemcpy3DParms myparams; + hipGraph_t graph; + hipGraphNode_t memcpyNode; + std::vector dependencies; + hipStream_t streamForGraph; + hipGraphExec_t graphExec; + + HIP_CHECK(hipSetDevice(0)); + HIP_CHECK(hipStreamCreate(&streamForGraph)); + // Initialize 1D object + for (int i = 0; i < XSIZE; i++) { + harray1D[i] = i + 1; + } + + hipChannelFormatDesc channelDesc = hipCreateChannelDesc(sizeof(int)*8, + 0, 0, 0, formatKind); + // Allocate 1D device array by passing depth(0), height(0) + HIP_CHECK(hipMalloc3DArray(&devArray1, &channelDesc, + make_hipExtent(width, 0, 0), hipArrayDefault)); + HIP_CHECK(hipMalloc3DArray(&devArray2, &channelDesc, + make_hipExtent(width, 0, 0), hipArrayDefault)); + HIP_CHECK(hipGraphCreate(&graph, 0)); + + // For peer access test, Memory is allocated on device(0) + // while memcpy nodes are allocated and assigned to peer device(1) + if (peerAccess) { + HIP_CHECK(hipSetDevice(1)); + } + + // Host to Device + memset(&myparams, 0x0, sizeof(hipMemcpy3DParms)); + myparams.srcPos = make_hipPos(0, 0, 0); + myparams.dstPos = make_hipPos(0, 0, 0); + myparams.extent = make_hipExtent(width, 1, 1); + myparams.srcPtr = make_hipPitchedPtr(harray1D, width * sizeof(int), + width, 1); + myparams.dstArray = devArray1; + myparams.kind = hipMemcpyHostToDevice; + + HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, nullptr, 0, &myparams)); + dependencies.push_back(memcpyNode); + + // Device to Device + memset(&myparams, 0x0, sizeof(hipMemcpy3DParms)); + myparams.srcPos = make_hipPos(0, 0, 0); + myparams.dstPos = make_hipPos(0, 0, 0); + myparams.srcArray = devArray1; + myparams.dstArray = devArray2; + myparams.extent = make_hipExtent(width, 1, 1); + myparams.kind = hipMemcpyDeviceToDevice; + + HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, dependencies.data(), + dependencies.size(), &myparams)); + dependencies.clear(); + dependencies.push_back(memcpyNode); + + // Device to host + memset(&myparams, 0x0, sizeof(hipMemcpy3DParms)); + myparams.srcPos = make_hipPos(0, 0, 0); + myparams.dstPos = make_hipPos(0, 0, 0); + myparams.extent = make_hipExtent(width, 1, 1); + myparams.dstPtr = make_hipPitchedPtr(harray1Dres, width * sizeof(int), + width, 1); + myparams.srcArray = devArray2; + myparams.kind = hipMemcpyDeviceToHost; + + HIP_CHECK(hipGraphAddMemcpyNode(&memcpyNode, graph, dependencies.data(), + dependencies.size(), &myparams)); + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Validate result + for (int i = 0; i < XSIZE; i++) { + if (harray1D[i] != harray1Dres[i]) { + INFO("harray1D: " << harray1D[i] << " harray1Dres: " << harray1Dres[i] + << " mismatch at : " << i); + REQUIRE(false); + } + } + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + hipFreeArray(devArray1); + hipFreeArray(devArray2); +} + +/** + * Basic Functional Tests adds memcpy nodes of types H2D, D2D and D2H to graph + * and verifies execution sequence by launching graph on default device. + * Tests also verify memcpy node addition with 1D, 2D and 3D objects. + */ +TEST_CASE("Unit_hipGraphAddMemcpyNode_BasicFunctional") { + SECTION("Memcpy with 3D array on default device") { + validateMemcpyNode3DArray(); + } + + SECTION("Memcpy with 2D array on default device") { + validateMemcpyNode2DArray(); + } + + SECTION("Memcpy with 1D array on default device") { + validateMemcpyNode1DArray(); + } +} + +/** + * Peer access tests adds and assigns memcpy nodes of types H2D, D2D and D2H + * to peer device. Memory allocations happen on device(0) and memcpy operations + * are performed from device(1). + * Tests also verify memcpy node addition with 1D, 2D and 3D objects. + */ +TEST_CASE("Unit_hipGraphAddMemcpyNode_PeerAccessFunctional") { + int numDevices{}, peerAccess{}; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + if (numDevices > 1) { + HIP_CHECK(hipDeviceCanAccessPeer(&peerAccess, 1, 0)); + } + + if (!peerAccess) { + WARN("Skipping test as peer device access is not found!"); + return; + } + + SECTION("Memcpy with 3D array on peer device") { + validateMemcpyNode3DArray(true); + } + + SECTION("Memcpy with 2D array on peer device") { + validateMemcpyNode2DArray(true); + } + + SECTION("Memcpy with 1D array on peer device") { + validateMemcpyNode1DArray(true); + } +}