From c90d83366f7dbe78c9111588576441eb2a1bcaba Mon Sep 17 00:00:00 2001 From: Rambabu Swargam Date: Fri, 6 Dec 2024 18:19:40 +0530 Subject: [PATCH] SWDEV-491363 - [catch2][dtest] Add and Enable tests for Graph APIs Details : 1) Add tests for hipDrvGraphExecMemsetNodeSetParams 2) Enable commented tests for hipDrvGraphMemcpyNodeSetParams Change-Id: Ia4771e1dea022b5f4316b66c0bd585f0f941ab77 --- catch/include/memcpy3d_tests_common.hh | 3 - catch/unit/graph/CMakeLists.txt | 1 + .../hipDrvGraphExecMemsetNodeSetParams.cc | 228 ++++++++++++++++++ .../graph/hipDrvGraphMemcpyNodeSetParams.cc | 12 +- 4 files changed, 236 insertions(+), 8 deletions(-) create mode 100644 catch/unit/graph/hipDrvGraphExecMemsetNodeSetParams.cc diff --git a/catch/include/memcpy3d_tests_common.hh b/catch/include/memcpy3d_tests_common.hh index 1b8c136e27..845e43779c 100644 --- a/catch/include/memcpy3d_tests_common.hh +++ b/catch/include/memcpy3d_tests_common.hh @@ -740,8 +740,6 @@ static inline bool operator==(const HIP_MEMCPY3D& lhs, const HIP_MEMCPY3D& rhs) return pos_eq && extent_eq && mem_eq; } -// APIs hipDrvGraphMemcpyNodeGetParams, hipDrvGraphMemcpyNodeSetParams are yet to be implemented in HIP runtime. -#if 0 template hipError_t DrvMemcpy3DGraphWrapper(DrvPtrVariant dst_ptr, hipPos dst_pos, DrvPtrVariant src_ptr, hipPos src_pos, hipExtent extent, hipMemcpyKind kind, @@ -774,7 +772,6 @@ hipError_t DrvMemcpy3DGraphWrapper(DrvPtrVariant dst_ptr, hipPos dst_pos, DrvPtr return hipSuccess; } -#endif //if 0 template hipError_t DrvMemcpy3DWrapper(DrvPtrVariant dst_ptr, hipPos dst_pos, DrvPtrVariant src_ptr, diff --git a/catch/unit/graph/CMakeLists.txt b/catch/unit/graph/CMakeLists.txt index c9cc8f3c38..6a69414168 100644 --- a/catch/unit/graph/CMakeLists.txt +++ b/catch/unit/graph/CMakeLists.txt @@ -160,6 +160,7 @@ set(TEST_SRC hipDeviceSetGraphMemAttribute.cc hipDeviceGetGraphMemAttribute.cc hipDeviceGraphMemTrim.cc + hipDrvGraphExecMemsetNodeSetParams.cc ) if(HIP_PLATFORM MATCHES "amd") diff --git a/catch/unit/graph/hipDrvGraphExecMemsetNodeSetParams.cc b/catch/unit/graph/hipDrvGraphExecMemsetNodeSetParams.cc new file mode 100644 index 0000000000..416e7a5dbc --- /dev/null +++ b/catch/unit/graph/hipDrvGraphExecMemsetNodeSetParams.cc @@ -0,0 +1,228 @@ +/* +Copyright (c) 2024 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. +*/ + +#include +#include + +const size_t width = 1024; +const size_t height = 1; +const int N = width * height; +const int value = 120; + +/** + * Test Description + * ------------------------ + * - Test case to validate basic functionality of + * hipDrvGraphExecMemsetNodeSetParams with below scenario, + * 1) Allocate device memory, host memory + * 2) Create Graph + * 3) Add Memset node to Graph to set the elements value to 120 + * 4) Add Memcpy node to Graph to copy from device to host memory + * 4) Take GraphExec and Instantiate the graph + * 5) Launch the Graph using GraphExec + * 6) Validate the host array and all elements should contain value 120 + * 7) Prapare new Memset node parameters with value to 121 + * 8) Update the Node parameter in the memset node with new Memset node + * parameters using hipDrvGraphExecMemsetNodeSetParams + * 9) Launch the Graph using GraphExec + * 10) Again Validate the host array and all elements should contain + * value 121 + * Test source + * ------------------------ + * - catch/unit/memory/hipDrvGraphExecMemsetNodeSetParams.cc + */ +TEST_CASE("Unit_hipDrvGraphExecMemsetNodeSetParams_BasicPositive") { + CHECK_IMAGE_SUPPORT + + HIP_CHECK(hipInit(0)); + hipDevice_t device; + hipCtx_t context; + HIP_CHECK(hipDeviceGet(&device, 0)); + HIP_CHECK(hipCtxCreate(&context, 0, device)); + + size_t pitch; + hipDeviceptr_t devMemSrc; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&devMemSrc), &pitch, + width, height)); + + char *hostMemDst = new char[N]; + REQUIRE(hostMemDst != nullptr); + for (int i = 0; i < N; i++) { + hostMemDst[i] = 0; + } + + hipGraph_t graph = nullptr; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + hipGraphNode_t memsetNode, memcpyNode; + + // Prepare memset node + HIP_MEMSET_NODE_PARAMS initialMemsetParams{}; + initialMemsetParams.dst = devMemSrc; + initialMemsetParams.pitch = pitch; + initialMemsetParams.elementSize = sizeof(char); + initialMemsetParams.width = width; + initialMemsetParams.height = height; + initialMemsetParams.value = value; + + HIP_CHECK(hipDrvGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, + &initialMemsetParams, context)); + + // Prepare memcpyNode + ::std::vector memcpyNodeDependencies; + memcpyNodeDependencies.push_back(memsetNode); + + HIP_MEMCPY3D memcpyParams{}; + memcpyParams.srcMemoryType = hipMemoryTypeDevice; + memcpyParams.srcDevice = devMemSrc; + memcpyParams.srcPitch = pitch; + memcpyParams.dstMemoryType = hipMemoryTypeHost; + memcpyParams.dstHost = hostMemDst; + memcpyParams.dstPitch = width; + memcpyParams.srcXInBytes = 0; + memcpyParams.srcY = 0; + memcpyParams.srcZ = 0; + memcpyParams.dstXInBytes = 0; + memcpyParams.dstY = 0; + memcpyParams.dstZ = 0; + memcpyParams.WidthInBytes = width; + memcpyParams.Height = height; + memcpyParams.Depth = 1; + + HIP_CHECK(hipDrvGraphAddMemcpyNode( + &memcpyNode, graph, memcpyNodeDependencies.data(), + memcpyNodeDependencies.size(), &memcpyParams, context)); + + hipGraphExec_t graphExec; + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, nullptr)); + HIP_CHECK(hipStreamSynchronize(0)); + + for (int i = 0; i < N; i++) { + REQUIRE(hostMemDst[i] == value); + } + + HIP_MEMSET_NODE_PARAMS newMemsetParams{}; + newMemsetParams.dst = devMemSrc; + newMemsetParams.pitch = pitch; + newMemsetParams.elementSize = sizeof(char); + newMemsetParams.width = width; + newMemsetParams.height = height; + newMemsetParams.value = value + 1; + + // Update with new memset node params + HIP_CHECK(hipDrvGraphExecMemsetNodeSetParams(graphExec, memsetNode, + &newMemsetParams, context)); + + for (int i = 0; i < N; i++) { + hostMemDst[i] = 0; + } + + HIP_CHECK(hipGraphLaunch(graphExec, nullptr)); + HIP_CHECK(hipStreamSynchronize(0)); + + for (int i = 0; i < N; i++) { + REQUIRE(hostMemDst[i] == (value + 1)); + } + + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipFree(reinterpret_cast(devMemSrc))); + delete[] hostMemDst; + HIP_CHECK(hipCtxDestroy(context)); +} + +/** + * Test Description + * ------------------------ + * - Test case to validate below scenarios, + * 1) With Invalid Graph Exec + * 2) With Invalid Node + * 3) With Invalid Node params + * 4) With Invalid destination address in Node params + * ------------------------ + * - catch/unit/memory/hipDrvGraphExecMemsetNodeSetParams.cc + */ +TEST_CASE("Unit_hipDrvGraphExecMemsetNodeSetParams_Negative") { + CHECK_IMAGE_SUPPORT + + HIP_CHECK(hipInit(0)); + hipDevice_t device; + hipCtx_t context; + HIP_CHECK(hipDeviceGet(&device, 0)); + HIP_CHECK(hipCtxCreate(&context, 0, device)); + + hipGraph_t graph = nullptr; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + size_t pitch; + hipDeviceptr_t devMemSrc; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&devMemSrc), &pitch, + width, height)); + + // Prepare memset node + HIP_MEMSET_NODE_PARAMS memsetParams{}; + memsetParams.dst = devMemSrc; + memsetParams.pitch = pitch; + memsetParams.elementSize = sizeof(char); + memsetParams.width = width; + memsetParams.height = height; + memsetParams.value = value; + + hipGraphNode_t memsetNode; + HIP_CHECK(hipDrvGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, + &memsetParams, context)); + + hipGraphExec_t graphExec; + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + + SECTION("Invalid Graph Exec") { + HIP_CHECK_ERROR(hipDrvGraphExecMemsetNodeSetParams(nullptr, memsetNode, + &memsetParams, context), + hipErrorInvalidValue); + } + SECTION("Invalid Node") { + HIP_CHECK_ERROR(hipDrvGraphExecMemsetNodeSetParams(graphExec, nullptr, + &memsetParams, context), + hipErrorInvalidValue); + } + SECTION("Invalid Node params") { + HIP_CHECK_ERROR(hipDrvGraphExecMemsetNodeSetParams(graphExec, memsetNode, + nullptr, context), + hipErrorInvalidValue); + } + SECTION("Invalid destination address in Node params") { + HIP_MEMSET_NODE_PARAMS memsetParams{}; + memsetParams.dst = 0; + memsetParams.pitch = pitch; + memsetParams.elementSize = sizeof(char); + memsetParams.width = width; + memsetParams.height = height; + memsetParams.value = value; + + HIP_CHECK_ERROR(hipDrvGraphExecMemsetNodeSetParams(graphExec, memsetNode, + nullptr, context), + hipErrorInvalidValue); + } + + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipFree(reinterpret_cast(devMemSrc))); + HIP_CHECK(hipCtxDestroy(context)); +} diff --git a/catch/unit/graph/hipDrvGraphMemcpyNodeSetParams.cc b/catch/unit/graph/hipDrvGraphMemcpyNodeSetParams.cc index a25c7508cf..4cf8c826d6 100644 --- a/catch/unit/graph/hipDrvGraphMemcpyNodeSetParams.cc +++ b/catch/unit/graph/hipDrvGraphMemcpyNodeSetParams.cc @@ -26,8 +26,6 @@ THE SOFTWARE. #include #include -// hipDrvGraphMemcpyNodeSetParams API is is yet to be implemented in HIP runtime. -#if 0 /** * @addtogroup hipDrvGraphMemcpyNodeSetParams hipDrvGraphMemcpyNodeSetParams * @{ @@ -85,6 +83,11 @@ TEST_CASE("Unit_hipDrvGraphMemcpyNodeSetParams_Positive_Basic") { Memcpy3DDeviceToDeviceShell( std::bind(DrvMemcpy3DGraphWrapper, _1, _2, _3, _4, _5, _6, context, _7)); } + /* + * Setting back the old context, since in the Memcpy3DDeviceToDeviceShell + * function the current context is getting modified with hipSetDevice calls + */ + HIP_CHECK(hipCtxSetCurrent(context)); } HIP_CHECK(hipCtxPopCurrent(&context)); @@ -188,7 +191,7 @@ TEST_CASE("Unit_hipDrvGraphMemcpyNodeSetParams_Negative_Parameters") { auto invalid_params = GetDrvMemcpy3DParms(invalid_ptr, dst_pos, src_ptr, src_pos, extent, kind); HIP_CHECK_ERROR(hipDrvGraphMemcpyNodeSetParams(node, &invalid_params), - hipErrorInvalidPitchValue); + hipErrorInvalidValue); } SECTION("srcPitch < width") { @@ -197,7 +200,7 @@ TEST_CASE("Unit_hipDrvGraphMemcpyNodeSetParams_Negative_Parameters") { auto invalid_params = GetDrvMemcpy3DParms(dst_ptr, dst_pos, invalid_ptr, src_pos, extent, kind); HIP_CHECK_ERROR(hipDrvGraphMemcpyNodeSetParams(node, &invalid_params), - hipErrorInvalidPitchValue); + hipErrorInvalidValue); } SECTION("dstPitch > max pitch") { @@ -314,7 +317,6 @@ TEST_CASE("Unit_hipDrvGraphMemcpyNodeSetParams_Negative_Parameters") { HIP_CHECK(hipCtxPopCurrent(&context)); HIP_CHECK(hipCtxDestroy(context)); } -#endif // if 0 /** * End doxygen group GraphTest.