diff --git a/catch/include/memcpy1d_tests_common.hh b/catch/include/memcpy1d_tests_common.hh index c14e6db444..d357d992cb 100644 --- a/catch/include/memcpy1d_tests_common.hh +++ b/catch/include/memcpy1d_tests_common.hh @@ -24,10 +24,10 @@ THE SOFTWARE. #include -#include #include -#include +#include #include +#include static inline unsigned int GenerateLinearAllocationFlagCombinations( const LinearAllocs allocation_type) { @@ -141,7 +141,6 @@ void MemcpyDeviceToDeviceShell(F memcpy_func, const hipStream_t kernel_stream = int can_access_peer = 0; HIP_CHECK(hipDeviceCanAccessPeer(&can_access_peer, src_device, dst_device)); if (!can_access_peer) { - INFO("Peer access cannot be enabled between devices " << src_device << " " << dst_device); return; } HIP_CHECK(hipDeviceEnablePeerAccess(dst_device, 0)); diff --git a/catch/unit/graph/CMakeLists.txt b/catch/unit/graph/CMakeLists.txt index e1ee3f1907..2d0a3c50ff 100644 --- a/catch/unit/graph/CMakeLists.txt +++ b/catch/unit/graph/CMakeLists.txt @@ -54,6 +54,7 @@ set(TEST_SRC hipGraphAddMemcpyNode1D.cc hipGraphAddChildGraphNode.cc hipGraphNodeGetType.cc + hipGraphExecMemcpyNodeSetParams1D_old.cc hipGraphExecMemcpyNodeSetParams1D.cc hipGraphGetEdges.cc hipGraphGetEdges_old.cc @@ -72,6 +73,8 @@ set(TEST_SRC hipGraphEventWaitNodeGetEvent.cc hipGraphExecMemcpyNodeSetParams.cc hipStreamBeginCapture.cc + hipGraphAddMemcpyNode1D_old.cc + hipGraphAddMemcpyNode1D.cc hipStreamBeginCapture_old.cc hipStreamIsCapturing.cc hipStreamIsCapturing_old.cc) @@ -105,6 +108,7 @@ set(TEST_SRC hipGraphLaunch.cc hipGraphLaunch_old.cc hipGraphMemcpyNodeSetParams1D.cc + hipGraphMemcpyNodeSetParams1D_old.cc hipGraphExecMemcpyNodeSetParamsToSymbol_old.cc hipGraphExecMemcpyNodeSetParamsToSymbol.cc hipGraphNodeGetDependentNodes.cc diff --git a/catch/unit/graph/hipGraphAddMemcpyNode1D.cc b/catch/unit/graph/hipGraphAddMemcpyNode1D.cc index 40c8ef1847..83ae815019 100644 --- a/catch/unit/graph/hipGraphAddMemcpyNode1D.cc +++ b/catch/unit/graph/hipGraphAddMemcpyNode1D.cc @@ -6,237 +6,179 @@ 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 + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY 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 +FITNESS 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 +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/** -Testcase Scenarios : -Functional - -1) Add 1D 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) 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. -3) Passing pDependencies as nullptr, api should return success. -4) When numDependencies is max(size_t) and pDependencies is not valid ptr, api expected to return error code. -5) When pDependencies is nullptr, but numDependencies is non-zero, api expected to return error. -6) When destination ptr is nullptr, api expected to return error code. -7) When source ptr is nullptr, api expected to return error code. -8) If count is more than allocated size for source and destination ptr, error code is returned. -9) If count is less than or equal to allocated size of source and destination ptr, api should return success. -*/ +#include #include -#include -#include -#include - -static void validateMemcpyNode1DArray(bool peerAccess) { - constexpr int SIZE{32}; - int harray1D[SIZE]{}; - int harray1Dres[SIZE]{}; - hipGraph_t graph; - hipArray_t devArray1, devArray2; - hipGraphNode_t memcpyH2D, memcpyD2H, memcpyD2D; - constexpr int numBytes{SIZE * sizeof(int)}; - hipStream_t streamForGraph; - hipGraphExec_t graphExec; - - HIP_CHECK(hipSetDevice(0)); - HIP_CHECK(hipStreamCreate(&streamForGraph)); - HIP_CHECK(hipMalloc(&devArray1, numBytes)); - HIP_CHECK(hipMalloc(&devArray2, numBytes)); - - // Initialize 1D object - for (int i = 0; i < SIZE; i++) { - harray1D[i] = i + 1; - } - - 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 (harray1D -> devArray1) - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D, graph, nullptr, 0, - devArray1, harray1D, numBytes, hipMemcpyHostToDevice)); - - // Device to Device (devArray1 -> devArray2) - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2D, graph, &memcpyH2D, 1, - devArray2, devArray1, numBytes, hipMemcpyDeviceToDevice)); - - // Device to host (devArray2 -> harray1Dres) - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H, graph, &memcpyD2D, 1, - harray1Dres, devArray2, numBytes, hipMemcpyDeviceToHost)); - - // 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 < SIZE; 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)); - HIP_CHECK(hipFree(devArray1)); - HIP_CHECK(hipFree(devArray2)); -} +#include +#include +#include "graph_tests_common.hh" /** - * Functional Tests adds memcpy 1D nodes of types H2D, D2D and D2H to graph - * and verifies execution sequence by launching graph. - * - * For Default device test: Memory allocations and memory operations - * are performed from device(0). - * For Peer device test: Memory allocations happen on device(0) and memcpy operations - * are performed from device(1). + * @addtogroup hipGraphAddMemcpyNode1D hipGraphAddMemcpyNode1D + * @{ + * @ingroup GraphTest + * `hipGraphAddMemcpyNode1D(hipGraphNode_t *pGraphNode, hipGraph_t graph, const hipGraphNode_t + * *pDependencies, size_t numDependencies, void *dst, const void *src, size_t count, hipMemcpyKind + * kind)` - Creates a 1D memcpy node and adds it to a graph */ -TEST_CASE("Unit_hipGraphAddMemcpyNode1D_Functional") { - SECTION("Memcpy with 1D array on default device") { - validateMemcpyNode1DArray(false); - } - - SECTION("Memcpy with 1D array on peer device") { - 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; - } - validateMemcpyNode1DArray(true); - } -} - - /** - * Negative Test for API hipGraphAddMemcpyNode1D + * Test Description + * ------------------------ + * - Verify basic API behavior. A Memcpy1D node is created with parameters set according to the + * test run, after which the graph is run and the memcpy results are verified. + * The test is run for all possible memcpy directions, with both the corresponding memcpy + * kind and hipMemcpyDefault, as well as half page and full page allocation sizes. + * Test source + * ------------------------ + * - unit/graph/hipGraphAddMemcpyNode1D.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ -TEST_CASE("Unit_hipGraphAddMemcpyNode1D_Negative") { - constexpr size_t N = 1024; - constexpr size_t Nbytes = N * sizeof(int); - int *A_d, *A_h; - hipGraph_t graph; - hipGraphNode_t memcpyNode{}; - hipError_t ret; +TEST_CASE("Unit_hipGraphAddMemcpyNode1D_Positive_Basic") { + constexpr auto f = [](void* dst, void* src, size_t count, hipMemcpyKind direction) { + hipGraph_t graph = nullptr; + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipGraphNode_t node = nullptr; + HIP_CHECK(hipGraphAddMemcpyNode1D(&node, graph, nullptr, 0, dst, src, count, direction)); + hipGraphExec_t graph_exec = nullptr; + HIP_CHECK(hipGraphInstantiate(&graph_exec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graph_exec, hipStreamPerThread)); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); - HIP_CHECK(hipMalloc(&A_d, Nbytes)); - HIP_CHECK(hipMalloc(&A_h, Nbytes)); + HIP_CHECK(hipGraphExecDestroy(graph_exec)); + HIP_CHECK(hipGraphDestroy(graph)); + + return hipSuccess; + }; + +#if HT_NVIDIA + MemcpyWithDirectionCommonTests(f); +#else + using namespace std::placeholders; + + SECTION("Device to host") { + MemcpyDeviceToHostShell(std::bind(f, _1, _2, _3, hipMemcpyDeviceToHost)); + } + + SECTION("Device to host with default kind") { + MemcpyDeviceToHostShell(std::bind(f, _1, _2, _3, hipMemcpyDefault)); + } + + SECTION("Host to device") { + MemcpyHostToDeviceShell(std::bind(f, _1, _2, _3, hipMemcpyHostToDevice)); + } + + SECTION("Host to device with default kind") { + MemcpyHostToDeviceShell(std::bind(f, _1, _2, _3, hipMemcpyDefault)); + } + +// Disabled on AMD due to defect - EXSWHTEC-209 +#if 0 + SECTION("Host to host") { + MemcpyHostToHostShell(std::bind(f, _1, _2, _3, hipMemcpyHostToHost)); + } + + SECTION("Host to host with default kind") { + MemcpyHostToHostShell(std::bind(f, _1, _2, _3, hipMemcpyDefault)); + } +#endif + + SECTION("Device to device") { + SECTION("Peer access enabled") { + MemcpyDeviceToDeviceShell(std::bind(f, _1, _2, _3, hipMemcpyDeviceToDevice)); + } + SECTION("Peer access disabled") { + MemcpyDeviceToDeviceShell(std::bind(f, _1, _2, _3, hipMemcpyDeviceToDevice)); + } + } + + SECTION("Device to device with default kind") { + SECTION("Peer access enabled") { + MemcpyDeviceToDeviceShell(std::bind(f, _1, _2, _3, hipMemcpyDefault)); + } + SECTION("Peer access disabled") { + MemcpyDeviceToDeviceShell(std::bind(f, _1, _2, _3, hipMemcpyDefault)); + } + } +#endif +} + +/** + * Test Description + * ------------------------ + * - Verify API behaviour with invalid arguments: + * -# node is nullptr + * -# graph is nullptr + * -# pDependencies is nullptr when numDependencies is not zero + * -# A node in pDependencies originates from a different graph + * -# numDependencies is invalid + * -# A node is duplicated in pDependencies + * -# dst is nullptr + * -# src is nullptr + * -# kind is an invalid enum value + * -# count is zero + * -# count is larger than dst allocation size + * -# count is larger than src allocation size + * Test source + * ------------------------ + * - unit/graph/hipGraphAddMemcpyNode1D.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipGraphAddMemcpyNode1D_Negative_Parameters") { + using namespace std::placeholders; + hipGraph_t graph = nullptr; HIP_CHECK(hipGraphCreate(&graph, 0)); + hipGraphNode_t node = nullptr; + int src[2] = {}, dst[2] = {}; - SECTION("Pass pGraphNode as nullptr") { - ret = hipGraphAddMemcpyNode1D(nullptr, graph, - nullptr, 0, A_d, A_h, Nbytes, hipMemcpyHostToDevice); - REQUIRE(hipErrorInvalidValue == ret); + GraphAddNodeCommonNegativeTests( + std::bind(hipGraphAddMemcpyNode1D, _1, _2, _3, _4, dst, src, sizeof(dst), hipMemcpyDefault), + graph); + + MemcpyWithDirectionCommonNegativeTests( + std::bind(hipGraphAddMemcpyNode1D, &node, graph, nullptr, 0, _1, _2, _3, _4), dst, src, + sizeof(dst), hipMemcpyDefault); + +// Disabled on AMD due to defect - EXSWHTEC-211 +#if HT_NVIDIA + SECTION("count == 0") { + HIP_CHECK_ERROR( + hipGraphAddMemcpyNode1D(&node, graph, nullptr, 0, dst, src, 0, hipMemcpyDefault), + hipErrorInvalidValue); } - SECTION("Pass graph as nullptr") { - ret = hipGraphAddMemcpyNode1D(&memcpyNode, nullptr, - nullptr, 0, A_d, A_h, Nbytes, hipMemcpyHostToDevice); - REQUIRE(hipErrorInvalidValue == ret); +#endif + + SECTION("count larger than dst allocation size") { + LinearAllocGuard dev_dst(LinearAllocs::hipMalloc, sizeof(int)); + HIP_CHECK_ERROR(hipGraphAddMemcpyNode1D(&node, graph, nullptr, 0, dev_dst.ptr(), src, + sizeof(src), hipMemcpyDefault), + hipErrorInvalidValue); } - SECTION("Pass pDependencies as nullptr") { - ret = hipGraphAddMemcpyNode1D(&memcpyNode, graph, - nullptr, 0, A_d, A_h, Nbytes, hipMemcpyHostToDevice); - REQUIRE(hipSuccess == ret); + + SECTION("count larger than src allocation size") { + LinearAllocGuard dev_src(LinearAllocs::hipMalloc, sizeof(int)); + HIP_CHECK_ERROR(hipGraphAddMemcpyNode1D(&node, graph, nullptr, 0, dst, dev_src.ptr(), + sizeof(dst), hipMemcpyDefault), + hipErrorInvalidValue); } - SECTION("Pass numDependencies is max and pDependencies is not valid ptr") { - ret = hipGraphAddMemcpyNode1D(&memcpyNode, graph, - nullptr, INT_MAX, A_d, A_h, Nbytes, hipMemcpyHostToDevice); - REQUIRE(hipErrorInvalidValue == ret); - } - SECTION("Pass pDependencies as nullptr, but numDependencies is non-zero") { - ret = hipGraphAddMemcpyNode1D(&memcpyNode, graph, - nullptr, 9, A_d, A_h, Nbytes, hipMemcpyHostToDevice); - REQUIRE(hipErrorInvalidValue == ret); - } - SECTION("Pass destination ptr as nullptr") { - ret = hipGraphAddMemcpyNode1D(&memcpyNode, graph, - nullptr, 0, nullptr, A_h, Nbytes, hipMemcpyHostToDevice); - REQUIRE(hipErrorInvalidValue == ret); - } - SECTION("Pass source ptr as nullptr") { - ret = hipGraphAddMemcpyNode1D(&memcpyNode, graph, - nullptr, 0, A_d, nullptr, Nbytes, hipMemcpyHostToDevice); - REQUIRE(hipErrorInvalidValue == ret); - } - SECTION("Pass count as more than allocated size for source ptr") { - ret = hipGraphAddMemcpyNode1D(&memcpyNode, graph, - nullptr, 0, A_d, A_h, Nbytes+10, hipMemcpyHostToDevice); - REQUIRE(hipErrorInvalidValue == ret); - } - SECTION("Pass count as less than allocated size for destination ptr") { - ret = hipGraphAddMemcpyNode1D(&memcpyNode, graph, - nullptr, 0, A_d, A_h, Nbytes-10, hipMemcpyHostToDevice); - REQUIRE(hipSuccess == ret); - } - HIP_CHECK(hipFree(A_d)); - 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)}; - - // 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); - - hipGraph_t graph; - hipStream_t streamForGraph; - hipGraphExec_t graphExec; - hipGraphNode_t memcpyH2H; - HIP_CHECK(hipGraphCreate(&graph, 0)); - HIP_CHECK(hipStreamCreate(&streamForGraph)); - - // Host to Host - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2H, graph, nullptr, 0, - B_h.data(), A_h.data(), numBytes, hipMemcpyHostToHost)); - - // 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(std::equal(A_h.begin(), A_h.end(), B_h.begin(), B_h.end())); -} diff --git a/catch/unit/graph/hipGraphAddMemcpyNode1D_old.cc b/catch/unit/graph/hipGraphAddMemcpyNode1D_old.cc new file mode 100644 index 0000000000..40c8ef1847 --- /dev/null +++ b/catch/unit/graph/hipGraphAddMemcpyNode1D_old.cc @@ -0,0 +1,242 @@ +/* +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 +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) Add 1D 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) 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. +3) Passing pDependencies as nullptr, api should return success. +4) When numDependencies is max(size_t) and pDependencies is not valid ptr, api expected to return error code. +5) When pDependencies is nullptr, but numDependencies is non-zero, api expected to return error. +6) When destination ptr is nullptr, api expected to return error code. +7) When source ptr is nullptr, api expected to return error code. +8) If count is more than allocated size for source and destination ptr, error code is returned. +9) If count is less than or equal to allocated size of source and destination ptr, api should return success. +*/ + +#include +#include +#include +#include + +static void validateMemcpyNode1DArray(bool peerAccess) { + constexpr int SIZE{32}; + int harray1D[SIZE]{}; + int harray1Dres[SIZE]{}; + hipGraph_t graph; + hipArray_t devArray1, devArray2; + hipGraphNode_t memcpyH2D, memcpyD2H, memcpyD2D; + constexpr int numBytes{SIZE * sizeof(int)}; + hipStream_t streamForGraph; + hipGraphExec_t graphExec; + + HIP_CHECK(hipSetDevice(0)); + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipMalloc(&devArray1, numBytes)); + HIP_CHECK(hipMalloc(&devArray2, numBytes)); + + // Initialize 1D object + for (int i = 0; i < SIZE; i++) { + harray1D[i] = i + 1; + } + + 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 (harray1D -> devArray1) + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D, graph, nullptr, 0, + devArray1, harray1D, numBytes, hipMemcpyHostToDevice)); + + // Device to Device (devArray1 -> devArray2) + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2D, graph, &memcpyH2D, 1, + devArray2, devArray1, numBytes, hipMemcpyDeviceToDevice)); + + // Device to host (devArray2 -> harray1Dres) + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H, graph, &memcpyD2D, 1, + harray1Dres, devArray2, numBytes, hipMemcpyDeviceToHost)); + + // 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 < SIZE; 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)); + HIP_CHECK(hipFree(devArray1)); + HIP_CHECK(hipFree(devArray2)); +} + + +/** + * Functional Tests adds memcpy 1D nodes of types H2D, D2D and D2H to graph + * and verifies execution sequence by launching graph. + * + * For Default device test: Memory allocations and memory operations + * are performed from device(0). + * For Peer device test: Memory allocations happen on device(0) and memcpy operations + * are performed from device(1). + */ +TEST_CASE("Unit_hipGraphAddMemcpyNode1D_Functional") { + SECTION("Memcpy with 1D array on default device") { + validateMemcpyNode1DArray(false); + } + + SECTION("Memcpy with 1D array on peer device") { + 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; + } + validateMemcpyNode1DArray(true); + } +} + + + +/** + * Negative Test for API hipGraphAddMemcpyNode1D + */ +TEST_CASE("Unit_hipGraphAddMemcpyNode1D_Negative") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + int *A_d, *A_h; + hipGraph_t graph; + hipGraphNode_t memcpyNode{}; + hipError_t ret; + + HIP_CHECK(hipMalloc(&A_d, Nbytes)); + HIP_CHECK(hipMalloc(&A_h, Nbytes)); + HIP_CHECK(hipGraphCreate(&graph, 0)); + + SECTION("Pass pGraphNode as nullptr") { + ret = hipGraphAddMemcpyNode1D(nullptr, graph, + nullptr, 0, A_d, A_h, Nbytes, hipMemcpyHostToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass graph as nullptr") { + ret = hipGraphAddMemcpyNode1D(&memcpyNode, nullptr, + nullptr, 0, A_d, A_h, Nbytes, hipMemcpyHostToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass pDependencies as nullptr") { + ret = hipGraphAddMemcpyNode1D(&memcpyNode, graph, + nullptr, 0, A_d, A_h, Nbytes, hipMemcpyHostToDevice); + REQUIRE(hipSuccess == ret); + } + SECTION("Pass numDependencies is max and pDependencies is not valid ptr") { + ret = hipGraphAddMemcpyNode1D(&memcpyNode, graph, + nullptr, INT_MAX, A_d, A_h, Nbytes, hipMemcpyHostToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass pDependencies as nullptr, but numDependencies is non-zero") { + ret = hipGraphAddMemcpyNode1D(&memcpyNode, graph, + nullptr, 9, A_d, A_h, Nbytes, hipMemcpyHostToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass destination ptr as nullptr") { + ret = hipGraphAddMemcpyNode1D(&memcpyNode, graph, + nullptr, 0, nullptr, A_h, Nbytes, hipMemcpyHostToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass source ptr as nullptr") { + ret = hipGraphAddMemcpyNode1D(&memcpyNode, graph, + nullptr, 0, A_d, nullptr, Nbytes, hipMemcpyHostToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass count as more than allocated size for source ptr") { + ret = hipGraphAddMemcpyNode1D(&memcpyNode, graph, + nullptr, 0, A_d, A_h, Nbytes+10, hipMemcpyHostToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass count as less than allocated size for destination ptr") { + ret = hipGraphAddMemcpyNode1D(&memcpyNode, graph, + nullptr, 0, A_d, A_h, Nbytes-10, hipMemcpyHostToDevice); + REQUIRE(hipSuccess == ret); + } + HIP_CHECK(hipFree(A_d)); + 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)}; + + // 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); + + hipGraph_t graph; + hipStream_t streamForGraph; + hipGraphExec_t graphExec; + hipGraphNode_t memcpyH2H; + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipStreamCreate(&streamForGraph)); + + // Host to Host + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2H, graph, nullptr, 0, + B_h.data(), A_h.data(), numBytes, hipMemcpyHostToHost)); + + // 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(std::equal(A_h.begin(), A_h.end(), B_h.begin(), B_h.end())); +} diff --git a/catch/unit/graph/hipGraphExecMemcpyNodeSetParams1D.cc b/catch/unit/graph/hipGraphExecMemcpyNodeSetParams1D.cc index 03e97d32e6..9a01d6d0ae 100644 --- a/catch/unit/graph/hipGraphExecMemcpyNodeSetParams1D.cc +++ b/catch/unit/graph/hipGraphExecMemcpyNodeSetParams1D.cc @@ -6,8 +6,10 @@ 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 WARRANTY OF ANY KIND, EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE @@ -17,182 +19,235 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/* -Testcase Scenarios : -Functional- -1) Instantiate a graph with memcpy node, obtain executable graph and update the - node params with set exec api call. Make sure they are taking effect. -Negative- -1) Pass hGraphExec as nullptr and check if api returns error. -2) Pass GraphNode as nullptr and check if api returns error. -3) Pass destination ptr is nullptr, api expected to return error code. -4) Pass source ptr is nullptr, api expected to return error code. -5) Pass count as zero, api expected to return error code. -6) Pass same pointer as source ptr and destination ptr, api expected to return error code. -7) Pass overlap memory address as source ptr and destination ptr, api expected to return error code. -7) Pass overlap memory as source ptr and destination ptr where source ptr is ahead of destination ptr, api expected to return error code. -8) Pass overlap memory as source ptr and destination ptr where destination ptr is ahead of source ptr, api expected to return error code. -9) If count is more than allocated size for source and destination ptr, api should return error code. -10) If count is less than allocated size for source and destination ptr, api should return error code. -11) Change the hipMemcpyKind from H2D to D2H but allocate pointer memory for H2D, api should return error code. -*/ +#include #include -#include -#include +#include +#include -/* Test verifies hipGraphExecMemcpyNodeSetParams1D API Negative scenarios. +#include "graph_tests_common.hh" + +/** + * @addtogroup hipGraphExecMemcpyNodeSetParams1D hipGraphExecMemcpyNodeSetParams1D + * @{ + * @ingroup GraphTest + * `hipGraphExecMemcpyNodeSetParams1D(hipGraphExec_t hGraphExec, hipGraphNode_t node, void *dst, + * const void *src, size_t count, hipMemcpyKind kind)` - Sets the parameters for a memcpy node in + * the given graphExec to perform a 1-dimensional copy */ -TEST_CASE("Unit_hipGraphExecMemcpyNodeSetParams1D_Negative") { - constexpr size_t N = 1024; - constexpr size_t Nbytes = N * sizeof(int); - int *A_d; - HIP_CHECK(hipMalloc(&A_d, Nbytes)); - int *A_h = reinterpret_cast(malloc(Nbytes)); - REQUIRE(A_h != nullptr); - memset(A_h, 0, Nbytes); +/** + * Test Description + * ------------------------ + * - Verify that node parameters get updated correctly by creating a node with valid but + * incorrect parameters, and the setting them to the correct values in the executable graph. The + * executable graph is run and the results of the memcpy verified. The test is run for all possible + * memcpy directions, with both the corresponding memcpy kind and hipMemcpyDefault, as well as half + * page and full page allocation sizes. Test source + * ------------------------ + * - unit/graph/hipGraphExecMemcpyNodeSetParams1D.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipGraphExecMemcpyNodeSetParams1D_Positive_Basic") { + constexpr auto f = [](void* dst, void* src, size_t count, hipMemcpyKind direction) { + hipGraph_t graph = nullptr; + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipGraphNode_t node = nullptr; + const auto offset_src = reinterpret_cast(src) + 1; + const auto offset_dst = reinterpret_cast(dst) + 1; + HIP_CHECK(hipGraphAddMemcpyNode1D(&node, graph, nullptr, 0, offset_dst, offset_src, count - 1, + direction)); + hipGraphExec_t graph_exec = nullptr; + HIP_CHECK(hipGraphInstantiate(&graph_exec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphExecMemcpyNodeSetParams1D(graph_exec, node, dst, src, count, direction)); + HIP_CHECK(hipGraphLaunch(graph_exec, hipStreamPerThread)); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); - hipError_t ret; - hipGraphNode_t memcpyH2D; - hipGraph_t graph; - hipGraphExec_t graphExec; + HIP_CHECK(hipGraphExecDestroy(graph_exec)); + HIP_CHECK(hipGraphDestroy(graph)); + return hipSuccess; + }; + +#if HT_NVIDIA + MemcpyWithDirectionCommonTests(f); +#else + using namespace std::placeholders; + + SECTION("Device to host") { + MemcpyDeviceToHostShell(std::bind(f, _1, _2, _3, hipMemcpyDeviceToHost)); + } + + SECTION("Host to device") { + MemcpyHostToDeviceShell(std::bind(f, _1, _2, _3, hipMemcpyHostToDevice)); + } + + SECTION("Device to device") { + SECTION("Peer access enabled") { + MemcpyDeviceToDeviceShell(std::bind(f, _1, _2, _3, hipMemcpyDeviceToDevice)); + } + SECTION("Peer access disabled") { + MemcpyDeviceToDeviceShell(std::bind(f, _1, _2, _3, hipMemcpyDeviceToDevice)); + } + } + + SECTION("Device to device with default kind") { + SECTION("Peer access enabled") { + MemcpyDeviceToDeviceShell(std::bind(f, _1, _2, _3, hipMemcpyDefault)); + } + SECTION("Peer access disabled") { + MemcpyDeviceToDeviceShell(std::bind(f, _1, _2, _3, hipMemcpyDefault)); + } + } + +// Disabled on AMD due to defect - EXSWHTEC-209 +#if 0 + SECTION("Host to host") { + MemcpyHostToHostShell(std::bind(f, _1, _2, _3, hipMemcpyHostToHost)); + } + + SECTION("Host to host with default kind") { + MemcpyHostToHostShell(std::bind(f, _1, _2, _3, hipMemcpyDefault)); + } +#endif + +// Disabled on AMD due to defect - EXSWHTEC-210 +#if 0 + SECTION("Device to host with default kind") { + MemcpyDeviceToHostShell(std::bind(f, _1, _2, _3, hipMemcpyDefault)); + } + + SECTION("Host to device with default kind") { + MemcpyHostToDeviceShell(std::bind(f, _1, _2, _3, hipMemcpyDefault)); + } +#endif + +#endif +} + +/** + * Test Description + * ------------------------ + * - Verify API behaviour with invalid arguments: + * -# pGraphExec is nullptr + * -# node is nullptr + * -# graph is nullptr + * -# pDependencies is nullptr when numDependencies is not zero + * -# A node in pDependencies originates from a different graph + * -# numDependencies is invalid + * -# A node is duplicated in pDependencies + * -# dst is nullptr + * -# src is nullptr + * -# kind is an invalid enum value + * -# count is zero + * -# count is larger than dst allocation size + * -# count is larger than src allocation size + * Test source + * ------------------------ + * - unit/graph/hipGraphAddMemcpyNode1D.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipGraphExecMemcpyNodeSetParams1D_Negative_Parameters") { + using namespace std::placeholders; + hipGraph_t graph = nullptr; HIP_CHECK(hipGraphCreate(&graph, 0)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D, graph, nullptr, 0, A_d, A_h, - Nbytes, hipMemcpyHostToDevice)); - // Instantiate the graph - HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0)); - SECTION("Pass hGraphExec as nullptr") { - ret = hipGraphExecMemcpyNodeSetParams1D(nullptr, memcpyH2D, A_d, A_h, - Nbytes, hipMemcpyHostToDevice); - REQUIRE(hipErrorInvalidValue == ret); + int src[2] = {}, dst[2] = {}; + + hipGraphNode_t node = nullptr; + HIP_CHECK( + hipGraphAddMemcpyNode1D(&node, graph, nullptr, 0, dst, src, sizeof(dst), hipMemcpyDefault)); + + hipGraphExec_t graph_exec = nullptr; + HIP_CHECK(hipGraphInstantiate(&graph_exec, graph, nullptr, nullptr, 0)); + + SECTION("pGraphExec == nullptr") { + HIP_CHECK_ERROR( + hipGraphExecMemcpyNodeSetParams1D(nullptr, node, dst, src, sizeof(dst), hipMemcpyDefault), + hipErrorInvalidValue); } - SECTION("Pass GraphNode as nullptr") { - ret = hipGraphExecMemcpyNodeSetParams1D(graphExec, nullptr, A_d, A_h, - Nbytes, hipMemcpyHostToDevice); - REQUIRE(hipErrorInvalidValue == ret); + + SECTION("node == nullptr") { + HIP_CHECK_ERROR(hipGraphExecMemcpyNodeSetParams1D(graph_exec, nullptr, dst, src, sizeof(dst), + hipMemcpyDefault), + hipErrorInvalidValue); } - SECTION("Pass destination ptr is nullptr") { - ret = hipGraphExecMemcpyNodeSetParams1D(graphExec, memcpyH2D, nullptr, A_h, - Nbytes, hipMemcpyHostToDevice); - REQUIRE(hipErrorInvalidValue == ret); + + MemcpyWithDirectionCommonNegativeTests( + std::bind(hipGraphExecMemcpyNodeSetParams1D, graph_exec, node, _1, _2, _3, _4), dst, src, + sizeof(dst), hipMemcpyDefault); + + SECTION("count == 0") { + HIP_CHECK_ERROR( + hipGraphExecMemcpyNodeSetParams1D(graph_exec, node, dst, src, 0, hipMemcpyDefault), + hipErrorInvalidValue); } - SECTION("Pass source ptr is nullptr") { - ret = hipGraphExecMemcpyNodeSetParams1D(graphExec, memcpyH2D, A_d, nullptr, - Nbytes, hipMemcpyHostToDevice); - REQUIRE(hipErrorInvalidValue == ret); + + SECTION("count larger than dst allocation size") { + LinearAllocGuard dev_dst(LinearAllocs::hipMalloc, sizeof(int)); + HIP_CHECK_ERROR(hipGraphExecMemcpyNodeSetParams1D(graph_exec, node, dev_dst.ptr(), src, + sizeof(src), hipMemcpyDefault), + hipErrorInvalidValue); } - SECTION("Pass count as zero") { - ret = hipGraphExecMemcpyNodeSetParams1D(graphExec, memcpyH2D, A_d, A_h, - 0, hipMemcpyHostToDevice); - REQUIRE(hipErrorInvalidValue == ret); + + SECTION("count larger than src allocation size") { + LinearAllocGuard dev_src(LinearAllocs::hipMalloc, sizeof(int)); + HIP_CHECK_ERROR(hipGraphExecMemcpyNodeSetParams1D(graph_exec, node, dst, dev_src.ptr(), + sizeof(dst), hipMemcpyDefault), + hipErrorInvalidValue); } - SECTION("Pass same pointer as source ptr and destination ptr") { - ret = hipGraphExecMemcpyNodeSetParams1D(graphExec, memcpyH2D, A_d, A_d, - Nbytes, hipMemcpyDeviceToDevice); - REQUIRE(hipErrorInvalidValue == ret); - } - SECTION("Pass overlap memory where destination ptr is ahead of source ptr") { - ret = hipGraphExecMemcpyNodeSetParams1D(graphExec, memcpyH2D, A_d, A_d-5, - Nbytes, hipMemcpyDeviceToDevice); - REQUIRE(hipErrorInvalidValue == ret); - } - SECTION("Pass overlap memory where source ptr is ahead of destination ptr") { - ret = hipGraphExecMemcpyNodeSetParams1D(graphExec, memcpyH2D, A_d+5, A_d, - Nbytes, hipMemcpyDeviceToDevice); - REQUIRE(hipErrorInvalidValue == ret); - } - SECTION("Copy more than allocated memory") { - ret = hipGraphExecMemcpyNodeSetParams1D(graphExec, memcpyH2D, A_d, A_h, - Nbytes+8, hipMemcpyHostToDevice); - REQUIRE(hipErrorInvalidValue == ret); - } - SECTION("Copy less than allocated memory") { - ret = hipGraphExecMemcpyNodeSetParams1D(graphExec, memcpyH2D, A_d, A_h, - Nbytes-8, hipMemcpyHostToDevice); - REQUIRE(hipSuccess == ret); - } - SECTION("Change the hipMemcpyKind from H2D to D2H") { - ret = hipGraphExecMemcpyNodeSetParams1D(graphExec, memcpyH2D, A_d, A_h, - Nbytes, hipMemcpyDeviceToHost); - REQUIRE(hipSuccess != ret); - } - HIP_CHECK(hipFree(A_d)); - free(A_h); - HIP_CHECK(hipGraphExecDestroy(graphExec)); + + HIP_CHECK(hipGraphExecDestroy(graph_exec)); HIP_CHECK(hipGraphDestroy(graph)); } -/* Test verifies hipGraphExecMemcpyNodeSetParams1D API Functional scenarios. +/** + * Test Description + * ------------------------ + * - Verify that memcpy direction cannot be altered in an executable graph. The test is run for + * all memcpy directions with appropriate memory allocations. + * Test source + * ------------------------ + * - unit/graph/hipGraphExecMemcpyNodeSetParams1D.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ -TEST_CASE("Unit_hipGraphExecMemcpyNodeSetParams1D_Functional") { - constexpr size_t N = 1024; - constexpr size_t Nbytes = N * sizeof(int); - constexpr auto blocksPerCU = 6; // to hide latency - constexpr auto threadsPerBlock = 256; - int *A_d, *B_d, *C_d; - int *A_h, *B_h, *C_h; - size_t NElem{N}; +TEST_CASE("Unit_hipGraphExecMemcpyNodeSetParams1D_Negative_Changing_Memcpy_Direction") { + int host; + LinearAllocGuard dev(LinearAllocs::hipMalloc, sizeof(int)); - int *hData = reinterpret_cast(malloc(Nbytes)); - REQUIRE(hData != nullptr); - memset(hData, 0, Nbytes); - - hipGraphNode_t memcpyH2D_A, memcpyH2D_B, memcpyD2H_C; - hipGraphNode_t kernel_vecAdd; - hipKernelNodeParams kernelNodeParams{}; - hipGraph_t graph; - hipGraphExec_t graphExec; - hipStream_t streamForGraph; - - HIP_CHECK(hipStreamCreate(&streamForGraph)); - - HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); - unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + const auto [dir, src, dst] = + GENERATE_REF(std::make_tuple(hipMemcpyHostToHost, &host, &host), + std::make_tuple(hipMemcpyHostToDevice, &host, dev.ptr()), + std::make_tuple(hipMemcpyDeviceToHost, dev.ptr(), &host), + std::make_tuple(hipMemcpyDeviceToDevice, dev.ptr(), dev.ptr())); + hipGraph_t graph = nullptr; HIP_CHECK(hipGraphCreate(&graph, 0)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, - Nbytes, hipMemcpyHostToDevice)); + hipGraphNode_t node = nullptr; + HIP_CHECK(hipGraphAddMemcpyNode1D(&node, graph, nullptr, 0, dst, src, sizeof(int), dir)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B, graph, nullptr, 0, B_d, B_h, - Nbytes, hipMemcpyHostToDevice)); + hipGraphExec_t graph_exec = nullptr; + HIP_CHECK(hipGraphInstantiate(&graph_exec, graph, nullptr, nullptr, 0)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_C, graph, nullptr, 0, C_h, C_d, - Nbytes, hipMemcpyDeviceToHost)); + const auto set_dir = GENERATE(hipMemcpyHostToHost, hipMemcpyHostToDevice, hipMemcpyDeviceToHost, + hipMemcpyDeviceToDevice, hipMemcpyDefault); + if (dir == set_dir) { + HIP_CHECK(hipGraphExecDestroy(graph_exec)); + HIP_CHECK(hipGraphDestroy(graph)); + return; + } - void* kernelArgs2[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; - kernelNodeParams.func = reinterpret_cast(HipTest::vectorADD); - kernelNodeParams.gridDim = dim3(blocks); - kernelNodeParams.blockDim = dim3(threadsPerBlock); - kernelNodeParams.sharedMemBytes = 0; - kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs2); - kernelNodeParams.extra = nullptr; - HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph, nullptr, 0, - &kernelNodeParams)); + HIP_CHECK_ERROR( + hipGraphExecMemcpyNodeSetParams1D(graph_exec, node, dst, src, sizeof(int), set_dir), + hipErrorInvalidValue); - // Create dependencies - HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, &kernel_vecAdd, 1)); - HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_B, &kernel_vecAdd, 1)); - HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecAdd, &memcpyD2H_C, 1)); - - // Instantiate the graph - HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); - - HIP_CHECK(hipGraphExecMemcpyNodeSetParams1D(graphExec, memcpyD2H_C, hData, - C_d, Nbytes, hipMemcpyDeviceToHost)); - - HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); - HIP_CHECK(hipStreamSynchronize(streamForGraph)); - - // Verify graph execution result - HipTest::checkVectorADD(A_h, B_h, hData, N); - - HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); - HIP_CHECK(hipGraphExecDestroy(graphExec)); - HIP_CHECK(hipStreamDestroy(streamForGraph)); + HIP_CHECK(hipGraphExecDestroy(graph_exec)); HIP_CHECK(hipGraphDestroy(graph)); - free(hData); } diff --git a/catch/unit/graph/hipGraphExecMemcpyNodeSetParams1D_old.cc b/catch/unit/graph/hipGraphExecMemcpyNodeSetParams1D_old.cc new file mode 100644 index 0000000000..0a28e6f31e --- /dev/null +++ b/catch/unit/graph/hipGraphExecMemcpyNodeSetParams1D_old.cc @@ -0,0 +1,201 @@ +/* +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 WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS 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 IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/* +Testcase Scenarios : +Functional- +1) Instantiate a graph with memcpy node, obtain executable graph and update the + node params with set exec api call. Make sure they are taking effect. +Negative- +1) Pass hGraphExec as nullptr and check if api returns error. +2) Pass GraphNode as nullptr and check if api returns error. +3) Pass destination ptr is nullptr, api expected to return error code. +4) Pass source ptr is nullptr, api expected to return error code. +5) Pass count as zero, api expected to return error code. +6) Pass same pointer as source ptr and destination ptr, api expected to return error code. +7) Pass overlap memory address as source ptr and destination ptr, api expected to return error code. +7) Pass overlap memory as source ptr and destination ptr where source ptr is ahead of destination ptr, api expected to return error code. +8) Pass overlap memory as source ptr and destination ptr where destination ptr is ahead of source ptr, api expected to return error code. +9) If count is more than allocated size for source and destination ptr, api should return error code. +10) If count is less than allocated size for source and destination ptr, api should return error code. +11) Change the hipMemcpyKind from H2D to D2H but allocate pointer memory for H2D, api should return error code. +*/ + +#include +#include +#include +#include + +/* Test verifies hipGraphExecMemcpyNodeSetParams1D API Functional scenarios. + */ +TEST_CASE("Unit_hipGraphExecMemcpyNodeSetParams1D_Functional") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + int *A_d, *B_d, *C_d; + int *A_h, *B_h, *C_h; + size_t NElem{N}; + + int *hData = reinterpret_cast(malloc(Nbytes)); + REQUIRE(hData != nullptr); + memset(hData, 0, Nbytes); + + hipGraphNode_t memcpyH2D_A, memcpyH2D_B, memcpyD2H_C; + hipGraphNode_t kernel_vecAdd; + hipKernelNodeParams kernelNodeParams{}; + hipGraph_t graph; + hipGraphExec_t graphExec; + hipStream_t streamForGraph; + + HIP_CHECK(hipStreamCreate(&streamForGraph)); + + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, + Nbytes, hipMemcpyHostToDevice)); + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B, graph, nullptr, 0, B_d, B_h, + Nbytes, hipMemcpyHostToDevice)); + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_C, graph, nullptr, 0, C_h, C_d, + Nbytes, hipMemcpyDeviceToHost)); + + void* kernelArgs2[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + kernelNodeParams.func = reinterpret_cast(HipTest::vectorADD); + kernelNodeParams.gridDim = dim3(blocks); + kernelNodeParams.blockDim = dim3(threadsPerBlock); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs2); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph, nullptr, 0, + &kernelNodeParams)); + + // Create dependencies + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, &kernel_vecAdd, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_B, &kernel_vecAdd, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecAdd, &memcpyD2H_C, 1)); + + // Instantiate the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + + HIP_CHECK(hipGraphExecMemcpyNodeSetParams1D(graphExec, memcpyD2H_C, hData, + C_d, Nbytes, hipMemcpyDeviceToHost)); + + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Verify graph execution result + HipTest::checkVectorADD(A_h, B_h, hData, N); + + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + HIP_CHECK(hipGraphDestroy(graph)); + free(hData); +} + +/* Test verifies hipGraphExecMemcpyNodeSetParams1D API Negative scenarios. + */ +TEST_CASE("Unit_hipGraphExecMemcpyNodeSetParams1D_Negative") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + + LinearAllocGuard A_d(LinearAllocs::hipMalloc, Nbytes); + LinearAllocGuard A_h(LinearAllocs::malloc, Nbytes); + memset(A_h.ptr(), 0, Nbytes); + + hipGraph_t graph; + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipGraphNode_t memcpyH2D; + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D, graph, nullptr, 0, A_d.ptr(), A_h.ptr(), + Nbytes, hipMemcpyHostToDevice)); + // Instantiate the graph + hipGraphExec_t graphExec; + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, NULL, NULL, 0)); + + SECTION("Pass hGraphExec as nullptr") { + HIP_CHECK_ERROR(hipGraphExecMemcpyNodeSetParams1D(nullptr, memcpyH2D, A_d.ptr(), + A_h.ptr(), Nbytes, + hipMemcpyHostToDevice), + hipErrorInvalidValue); + } + + SECTION("Pass GraphNode as nullptr") { + HIP_CHECK_ERROR(hipGraphExecMemcpyNodeSetParams1D(graphExec, nullptr, A_d.ptr(), + A_h.ptr(), Nbytes, + hipMemcpyHostToDevice), + hipErrorInvalidValue); + } + + SECTION("Pass destination ptr is nullptr") { + HIP_CHECK_ERROR(hipGraphExecMemcpyNodeSetParams1D(graphExec, memcpyH2D, + nullptr, A_h.ptr(), Nbytes, + hipMemcpyHostToDevice), + hipErrorInvalidValue); + } + + SECTION("Pass source ptr is nullptr") { + HIP_CHECK_ERROR(hipGraphExecMemcpyNodeSetParams1D(graphExec, memcpyH2D, A_d.ptr(), + nullptr, Nbytes, + hipMemcpyHostToDevice), + hipErrorInvalidValue); + } + + SECTION("Pass count as zero") { + HIP_CHECK_ERROR(hipGraphExecMemcpyNodeSetParams1D(graphExec, memcpyH2D, A_d.ptr(), + A_h.ptr(), 0, + hipMemcpyHostToDevice), + hipErrorInvalidValue); + } + + SECTION("Pass same pointer as source ptr and destination ptr") { + HIP_CHECK_ERROR(hipGraphExecMemcpyNodeSetParams1D(graphExec, memcpyH2D, A_d.ptr(), + A_d.ptr(), Nbytes, + hipMemcpyDeviceToDevice), + hipErrorInvalidValue); + } + + SECTION("Pass overlap memory where destination ptr is ahead of source ptr") { + HIP_CHECK_ERROR(hipGraphExecMemcpyNodeSetParams1D(graphExec, memcpyH2D, A_d.ptr(), + A_d.ptr() - 5, Nbytes, + hipMemcpyDeviceToDevice), + hipErrorInvalidValue); + } + + SECTION("Pass overlap memory where source ptr is ahead of destination ptr") { + HIP_CHECK_ERROR(hipGraphExecMemcpyNodeSetParams1D(graphExec, memcpyH2D, + A_d.ptr() + 5, A_d.ptr(), Nbytes, + hipMemcpyDeviceToDevice), + hipErrorInvalidValue); + } + + SECTION("Copy more than allocated memory") { + HIP_CHECK_ERROR(hipGraphExecMemcpyNodeSetParams1D(graphExec, memcpyH2D, A_d.ptr(), + A_h.ptr(), Nbytes + 8, + hipMemcpyHostToDevice), + hipErrorInvalidValue); + } + + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); +} diff --git a/catch/unit/graph/hipGraphMemcpyNodeSetParams1D.cc b/catch/unit/graph/hipGraphMemcpyNodeSetParams1D.cc index 86e439e528..fa22c6ee2c 100644 --- a/catch/unit/graph/hipGraphMemcpyNodeSetParams1D.cc +++ b/catch/unit/graph/hipGraphMemcpyNodeSetParams1D.cc @@ -6,8 +6,10 @@ 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 WARRANTY OF ANY KIND, EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE @@ -17,169 +19,180 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/** -Testcase Scenarios : -Functional- -1) Create a graph, add Memcpy node to graph, update the Memcpy node params with set and make sure they are taking effect. -Negative- -1) Pass pGraphNode as nullptr and check if api returns error. -2) Pass destination ptr is nullptr, api expected to return error code. -3) Pass source ptr is nullptr, api expected to return error code. -4) Pass count as zero, api expected to return error code. -5) Pass same pointer as source ptr and destination ptr, api expected to return error code. -6) Pass overlap memory as source ptr and destination ptr where source ptr is ahead of destination ptr, api expected to return error code. -7) Pass overlap memory as source ptr and destination ptr where destination ptr is ahead of source ptr, api expected to return error code. -8) If count is more than allocated size for source and destination ptr, api should return error code. -9) If count is less than allocated size for source and destination ptr, api should return error code. -*/ +#include #include -#include -#include +#include +#include -/* Test verifies hipGraphMemcpyNodeSetParams1D API Negative scenarios. +#include "graph_tests_common.hh" + +static inline hipMemcpyKind ReverseMemcpyDirection(const hipMemcpyKind direction) { + switch (direction) { + case hipMemcpyHostToDevice: + return hipMemcpyDeviceToHost; + case hipMemcpyDeviceToHost: + return hipMemcpyHostToDevice; + default: + return direction; + } +}; + +/** + * @addtogroup hipGraphMemcpyNodeSetParams1D hipGraphMemcpyNodeSetParams1D + * @{ + * @ingroup GraphTest + * `hipGraphMemcpyNodeSetParams1D(hipGraphNode_t node, void *dst, const void *src, size_t count, + * hipMemcpyKind kind)` - Sets a memcpy node's parameters to perform a 1-dimensional copy */ -TEST_CASE("Unit_hipGraphMemcpyNodeSetParams1D_Negative") { - constexpr size_t N = 1024; - constexpr size_t Nbytes = N * sizeof(int); - int *A_d, *A_h; - hipGraphNode_t memcpyNode{}; - hipError_t ret; - HIP_CHECK(hipMalloc(&A_d, Nbytes)); - HIP_CHECK(hipMalloc(&A_h, Nbytes)); +/** + * Test Description + * ------------------------ + * - Verify that node parameters get updated correctly by creating a node with valid but + * incorrect parameters, and the setting them to the correct values after which the graph is + * executed and the results of the memcpy verified. + * The test is run for all possible memcpy directions, with both the corresponding memcpy + * kind and hipMemcpyDefault, as well as half page and full page allocation sizes. + * Test source + * ------------------------ + * - unit/graph/hipGraphMemcpyNodeSetParams1D.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 + */ +TEST_CASE("Unit_hipGraphMemcpyNodeSetParams1D_Positive_Basic") { + constexpr auto f = [](void* dst, void* src, size_t count, hipMemcpyKind direction) { + hipGraph_t graph = nullptr; + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipGraphNode_t node = nullptr; + HIP_CHECK(hipGraphAddMemcpyNode1D(&node, graph, nullptr, 0, src, dst, count / 2, + ReverseMemcpyDirection(direction))); + HIP_CHECK(hipGraphMemcpyNodeSetParams1D(node, dst, src, count, direction)); + hipGraphExec_t graph_exec = nullptr; + HIP_CHECK(hipGraphInstantiate(&graph_exec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graph_exec, hipStreamPerThread)); + HIP_CHECK(hipStreamSynchronize(hipStreamPerThread)); - hipGraph_t graph; - HIP_CHECK(hipGraphCreate(&graph, 0)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, A_d, A_h, - Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphExecDestroy(graph_exec)); + HIP_CHECK(hipGraphDestroy(graph)); - SECTION("Pass pGraphNode as nullptr") { - ret = hipGraphMemcpyNodeSetParams1D(nullptr, A_d, A_h, Nbytes, - hipMemcpyHostToDevice); - REQUIRE(hipErrorInvalidValue == ret); + return hipSuccess; + }; + +#if HT_NVIDIA + MemcpyWithDirectionCommonTests(f); +#else + using namespace std::placeholders; + + SECTION("Device to host") { + MemcpyDeviceToHostShell(std::bind(f, _1, _2, _3, hipMemcpyDeviceToHost)); } - SECTION("Pass destination ptr is nullptr") { - ret = hipGraphMemcpyNodeSetParams1D(memcpyNode, nullptr, A_h, Nbytes, - hipMemcpyHostToDevice); - REQUIRE(hipErrorInvalidValue == ret); + + SECTION("Host to device") { + MemcpyHostToDeviceShell(std::bind(f, _1, _2, _3, hipMemcpyHostToDevice)); } - SECTION("Pass source ptr is nullptr") { - ret = hipGraphMemcpyNodeSetParams1D(memcpyNode, A_d, nullptr, Nbytes, - hipMemcpyHostToDevice); - REQUIRE(hipErrorInvalidValue == ret); + + SECTION("Device to device") { + SECTION("Peer access enabled") { + MemcpyDeviceToDeviceShell(std::bind(f, _1, _2, _3, hipMemcpyDeviceToDevice)); + } + SECTION("Peer access disabled") { + MemcpyDeviceToDeviceShell(std::bind(f, _1, _2, _3, hipMemcpyDeviceToDevice)); + } } - SECTION("Pass count as zero") { - ret = hipGraphMemcpyNodeSetParams1D(memcpyNode, A_d, A_h, 0, - hipMemcpyHostToDevice); - REQUIRE(hipErrorInvalidValue == ret); + + SECTION("Device to device with default kind") { + SECTION("Peer access enabled") { + MemcpyDeviceToDeviceShell(std::bind(f, _1, _2, _3, hipMemcpyDefault)); + } + SECTION("Peer access disabled") { + MemcpyDeviceToDeviceShell(std::bind(f, _1, _2, _3, hipMemcpyDefault)); + } } -#if HT_AMD - SECTION("Pass same pointer as source ptr and destination ptr") { - ret = hipGraphMemcpyNodeSetParams1D(memcpyNode, A_d, A_d, Nbytes, - hipMemcpyDeviceToDevice); - REQUIRE(hipErrorInvalidValue == ret); + +// Disabled on AMD due to defect - EXSWHTEC-209 +#if 0 + SECTION("Host to host") { + MemcpyHostToHostShell(std::bind(f, _1, _2, _3, hipMemcpyHostToHost)); + } + + SECTION("Host to host with default kind") { + MemcpyHostToHostShell(std::bind(f, _1, _2, _3, hipMemcpyDefault)); } #endif - SECTION("Pass overlap memory where destination ptr is ahead of source ptr") { - ret = hipGraphMemcpyNodeSetParams1D(memcpyNode, A_d, A_d-5, Nbytes, - hipMemcpyDeviceToDevice); - REQUIRE(hipErrorInvalidValue == ret); - } - SECTION("Pass overlap memory where source ptr is ahead of destination ptr") { - ret = hipGraphMemcpyNodeSetParams1D(memcpyNode, A_d+5, A_d, Nbytes-5, - hipMemcpyDeviceToDevice); - REQUIRE(hipErrorInvalidValue == ret); - } - SECTION("Copy more than allocated memory") { - ret = hipGraphMemcpyNodeSetParams1D(memcpyNode, A_d, A_h, Nbytes+8, - hipMemcpyHostToDevice); - REQUIRE(hipErrorInvalidValue == ret); - } - SECTION("Copy less than allocated memory") { - ret = hipGraphMemcpyNodeSetParams1D(memcpyNode, A_d, A_h, Nbytes-8, - hipMemcpyHostToDevice); - REQUIRE(hipSuccess == ret); - } - SECTION("Change the kind from H2D to D2H") { - ret = hipGraphMemcpyNodeSetParams1D(memcpyNode, A_d, A_h, Nbytes, - hipMemcpyDeviceToHost); - REQUIRE(hipSuccess == ret); + +// Disabled on AMD due to defect - EXSWHTEC-210 +#if 0 + SECTION("Device to host with default kind") { + MemcpyDeviceToHostShell(std::bind(f, _1, _2, _3, hipMemcpyDefault)); } - HIP_CHECK(hipFree(A_d)); - HIP_CHECK(hipFree(A_h)); - HIP_CHECK(hipGraphDestroy(graph)); + SECTION("Host to device with default kind") { + MemcpyHostToDeviceShell(std::bind(f, _1, _2, _3, hipMemcpyDefault)); + } +#endif + +#endif } -/* Test verifies hipGraphMemcpyNodeSetParams1D API Functional scenarios. +/** + * Test Description + * ------------------------ + * - Verify API behaviour with invalid arguments: + * -# node is nullptr + * -# dst is nullptr + * -# src is nullptr + * -# kind is an invalid enum value + * -# count is zero + * -# count is larger than dst allocation size + * -# count is larger than src allocation size + * Test source + * ------------------------ + * - unit/graph/hipGraphMemcpyNodeSetParams1D.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 5.2 */ -TEST_CASE("Unit_hipGraphMemcpyNodeSetParams1D_Functional") { - constexpr size_t N = 1024; - constexpr size_t Nbytes = N * sizeof(int); - constexpr auto blocksPerCU = 6; // to hide latency - constexpr auto threadsPerBlock = 256; - int *A_d, *B_d, *C_d; - int *A_h, *B_h, *C_h; - size_t NElem{N}; - - int *hData = reinterpret_cast(malloc(Nbytes)); - REQUIRE(hData != nullptr); - memset(hData, 0, Nbytes); - - hipGraphNode_t memcpyH2D_A, memcpyH2D_B, memcpyD2H_C; - hipGraphNode_t kernel_vecAdd; - hipKernelNodeParams kernelNodeParams{}; - hipGraph_t graph; - hipGraphExec_t graphExec; - hipStream_t streamForGraph; - - HIP_CHECK(hipStreamCreate(&streamForGraph)); - - HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); - unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); - +TEST_CASE("Unit_hipGraphMemcpyNodeSetParams1D_Negative_Parameters") { + using namespace std::placeholders; + hipGraph_t graph = nullptr; HIP_CHECK(hipGraphCreate(&graph, 0)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, - Nbytes, hipMemcpyHostToDevice)); + int src[2] = {}, dst[2] = {}; - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B, graph, nullptr, 0, B_d, B_h, - Nbytes, hipMemcpyHostToDevice)); + hipGraphNode_t node = nullptr; + HIP_CHECK( + hipGraphAddMemcpyNode1D(&node, graph, nullptr, 0, dst, src, sizeof(dst), hipMemcpyDefault)); - HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_C, graph, nullptr, 0, C_h, C_d, - Nbytes, hipMemcpyDeviceToHost)); - HIP_CHECK(hipGraphMemcpyNodeSetParams1D(memcpyD2H_C, hData, C_d, Nbytes, - hipMemcpyDeviceToHost)); + SECTION("node == nullptr") { + HIP_CHECK_ERROR(hipGraphMemcpyNodeSetParams1D(nullptr, dst, src, sizeof(dst), hipMemcpyDefault), + hipErrorInvalidValue); + } - void* kernelArgs2[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; - kernelNodeParams.func = reinterpret_cast(HipTest::vectorADD); - kernelNodeParams.gridDim = dim3(blocks); - kernelNodeParams.blockDim = dim3(threadsPerBlock); - kernelNodeParams.sharedMemBytes = 0; - kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs2); - kernelNodeParams.extra = nullptr; - HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph, nullptr, 0, - &kernelNodeParams)); + MemcpyWithDirectionCommonNegativeTests( + std::bind(hipGraphMemcpyNodeSetParams1D, node, _1, _2, _3, _4), dst, src, sizeof(dst), + hipMemcpyDefault); - // Create dependencies - HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, &kernel_vecAdd, 1)); - HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_B, &kernel_vecAdd, 1)); - HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecAdd, &memcpyD2H_C, 1)); + SECTION("count == 0") { + HIP_CHECK_ERROR(hipGraphMemcpyNodeSetParams1D(node, dst, src, 0, hipMemcpyDefault), + hipErrorInvalidValue); + } - // Instantiate and launch the graph - HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); - HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); - HIP_CHECK(hipStreamSynchronize(streamForGraph)); + SECTION("count larger than dst allocation size") { + LinearAllocGuard dev_dst(LinearAllocs::hipMalloc, sizeof(int)); + HIP_CHECK_ERROR( + hipGraphMemcpyNodeSetParams1D(node, dev_dst.ptr(), src, sizeof(src), hipMemcpyDefault), + hipErrorInvalidValue); + } - // Verify graph execution result - HipTest::checkVectorADD(A_h, B_h, hData, N); + SECTION("count larger than src allocation size") { + LinearAllocGuard dev_src(LinearAllocs::hipMalloc, sizeof(int)); + HIP_CHECK_ERROR( + hipGraphMemcpyNodeSetParams1D(node, dst, dev_src.ptr(), sizeof(dst), hipMemcpyDefault), + hipErrorInvalidValue); + } - HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); - HIP_CHECK(hipGraphExecDestroy(graphExec)); - HIP_CHECK(hipStreamDestroy(streamForGraph)); HIP_CHECK(hipGraphDestroy(graph)); - free(hData); } - diff --git a/catch/unit/graph/hipGraphMemcpyNodeSetParams1D_old.cc b/catch/unit/graph/hipGraphMemcpyNodeSetParams1D_old.cc new file mode 100644 index 0000000000..414eda51e9 --- /dev/null +++ b/catch/unit/graph/hipGraphMemcpyNodeSetParams1D_old.cc @@ -0,0 +1,172 @@ +/* +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 WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS 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 IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/** +Testcase Scenarios : +Functional- +1) Create a graph, add Memcpy node to graph, update the Memcpy node params with set and make sure +they are taking effect. Negative- 1) Pass pGraphNode as nullptr and check if api returns error. 2) +Pass destination ptr is nullptr, api expected to return error code. 3) Pass source ptr is nullptr, +api expected to return error code. 4) Pass count as zero, api expected to return error code. 5) Pass +same pointer as source ptr and destination ptr, api expected to return error code. 6) Pass overlap +memory as source ptr and destination ptr where source ptr is ahead of destination ptr, api expected +to return error code. 7) Pass overlap memory as source ptr and destination ptr where destination ptr +is ahead of source ptr, api expected to return error code. 8) If count is more than allocated size +for source and destination ptr, api should return error code. 9) If count is less than allocated +size for source and destination ptr, api should return error code. +*/ + +#include +#include +#include + +/* Test verifies hipGraphMemcpyNodeSetParams1D API Negative scenarios. + */ +TEST_CASE("Unit_hipGraphMemcpyNodeSetParams1D_Negative") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + int *A_d, *A_h; + hipGraphNode_t memcpyNode{}; + hipError_t ret; + + HIP_CHECK(hipMalloc(&A_d, Nbytes)); + HIP_CHECK(hipMalloc(&A_h, Nbytes)); + + hipGraph_t graph; + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyNode, graph, nullptr, 0, A_d, A_h, Nbytes, + hipMemcpyHostToDevice)); + + SECTION("Pass pGraphNode as nullptr") { + ret = hipGraphMemcpyNodeSetParams1D(nullptr, A_d, A_h, Nbytes, hipMemcpyHostToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass destination ptr is nullptr") { + ret = hipGraphMemcpyNodeSetParams1D(memcpyNode, nullptr, A_h, Nbytes, hipMemcpyHostToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass source ptr is nullptr") { + ret = hipGraphMemcpyNodeSetParams1D(memcpyNode, A_d, nullptr, Nbytes, hipMemcpyHostToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass count as zero") { + ret = hipGraphMemcpyNodeSetParams1D(memcpyNode, A_d, A_h, 0, hipMemcpyHostToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } +#if HT_AMD + SECTION("Pass same pointer as source ptr and destination ptr") { + ret = hipGraphMemcpyNodeSetParams1D(memcpyNode, A_d, A_d, Nbytes, hipMemcpyDeviceToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } +#endif + SECTION("Pass overlap memory where destination ptr is ahead of source ptr") { + ret = hipGraphMemcpyNodeSetParams1D(memcpyNode, A_d, A_d - 5, Nbytes, hipMemcpyDeviceToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Pass overlap memory where source ptr is ahead of destination ptr") { + ret = hipGraphMemcpyNodeSetParams1D(memcpyNode, A_d + 5, A_d, Nbytes - 5, + hipMemcpyDeviceToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Copy more than allocated memory") { + ret = hipGraphMemcpyNodeSetParams1D(memcpyNode, A_d, A_h, Nbytes + 8, hipMemcpyHostToDevice); + REQUIRE(hipErrorInvalidValue == ret); + } + SECTION("Copy less than allocated memory") { + ret = hipGraphMemcpyNodeSetParams1D(memcpyNode, A_d, A_h, Nbytes - 8, hipMemcpyHostToDevice); + REQUIRE(hipSuccess == ret); + } + SECTION("Change the kind from H2D to D2H") { + ret = hipGraphMemcpyNodeSetParams1D(memcpyNode, A_d, A_h, Nbytes, hipMemcpyDeviceToHost); + REQUIRE(hipSuccess == ret); + } + + HIP_CHECK(hipFree(A_d)); + HIP_CHECK(hipFree(A_h)); + HIP_CHECK(hipGraphDestroy(graph)); +} + +/* Test verifies hipGraphMemcpyNodeSetParams1D API Functional scenarios. + */ +TEST_CASE("Unit_hipGraphMemcpyNodeSetParams1D_Functional") { + constexpr size_t N = 1024; + constexpr size_t Nbytes = N * sizeof(int); + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + int *A_d, *B_d, *C_d; + int *A_h, *B_h, *C_h; + size_t NElem{N}; + + int* hData = reinterpret_cast(malloc(Nbytes)); + REQUIRE(hData != nullptr); + memset(hData, 0, Nbytes); + + hipGraphNode_t memcpyH2D_A, memcpyH2D_B, memcpyD2H_C; + hipGraphNode_t kernel_vecAdd; + hipKernelNodeParams kernelNodeParams{}; + hipGraph_t graph; + hipGraphExec_t graphExec; + hipStream_t streamForGraph; + + HIP_CHECK(hipStreamCreate(&streamForGraph)); + + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N, false); + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_d, A_h, Nbytes, + hipMemcpyHostToDevice)); + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B, graph, nullptr, 0, B_d, B_h, Nbytes, + hipMemcpyHostToDevice)); + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyD2H_C, graph, nullptr, 0, C_h, C_d, Nbytes, + hipMemcpyDeviceToHost)); + + HIP_CHECK(hipGraphMemcpyNodeSetParams1D(memcpyD2H_C, hData, C_d, Nbytes, hipMemcpyDeviceToHost)); + + void* kernelArgs2[] = {&A_d, &B_d, &C_d, reinterpret_cast(&NElem)}; + kernelNodeParams.func = reinterpret_cast(HipTest::vectorADD); + kernelNodeParams.gridDim = dim3(blocks); + kernelNodeParams.blockDim = dim3(threadsPerBlock); + kernelNodeParams.sharedMemBytes = 0; + kernelNodeParams.kernelParams = reinterpret_cast(kernelArgs2); + kernelNodeParams.extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernel_vecAdd, graph, nullptr, 0, &kernelNodeParams)); + + // Create dependencies + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_A, &kernel_vecAdd, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_B, &kernel_vecAdd, 1)); + HIP_CHECK(hipGraphAddDependencies(graph, &kernel_vecAdd, &memcpyD2H_C, 1)); + + // Instantiate and launch the graph + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Verify graph execution result + HipTest::checkVectorADD(A_h, B_h, hData, N); + + HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + HIP_CHECK(hipGraphDestroy(graph)); + free(hData); +}