From ad4d04fa88f595afb324b40d19fa7af5baeb33d6 Mon Sep 17 00:00:00 2001 From: Nives Vukovic Date: Mon, 25 Dec 2023 17:57:36 +0000 Subject: [PATCH] EXSWHTEC-371 - Implement tests for the hipDrvGraphAddMemsetNode API #447 Change-Id: I981e14f34b008054d46e61ebae0099792df446b1 [ROCm/hip-tests commit: 63c93f95a89423d8fd8cfd20922c7bec55ffc5f5] --- .../catch/include/memcpy3d_tests_common.hh | 2 +- .../hip-tests/catch/unit/graph/CMakeLists.txt | 1 + .../graph/graph_memset_node_test_common.hh | 6 +- .../unit/graph/hipDrvGraphAddMemsetNode.cc | 670 ++++++++++++++++++ .../catch/unit/graph/hipGraphAddMemsetNode.cc | 2 +- .../unit/graph/hipGraphMemsetNodeSetParams.cc | 2 +- 6 files changed, 677 insertions(+), 6 deletions(-) create mode 100644 projects/hip-tests/catch/unit/graph/hipDrvGraphAddMemsetNode.cc diff --git a/projects/hip-tests/catch/include/memcpy3d_tests_common.hh b/projects/hip-tests/catch/include/memcpy3d_tests_common.hh index acb549c94a..68e388da19 100644 --- a/projects/hip-tests/catch/include/memcpy3d_tests_common.hh +++ b/projects/hip-tests/catch/include/memcpy3d_tests_common.hh @@ -881,4 +881,4 @@ void DrvMemcpy3DArrayDeviceShell(F memcpy_func, const hipStream_t kernel_stream }; PitchedMemoryVerify(host_alloc.ptr(), extent.width, extent.width / sizeof(int), extent.height, extent.depth, f); -} +} \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/graph/CMakeLists.txt b/projects/hip-tests/catch/unit/graph/CMakeLists.txt index 82525991c7..c7b4036c49 100644 --- a/projects/hip-tests/catch/unit/graph/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/graph/CMakeLists.txt @@ -153,6 +153,7 @@ set(TEST_SRC hipGraphAddMemFreeNode.cc hipDrvGraphMemcpyNodeGetParams.cc hipDrvGraphMemcpyNodeSetParams.cc + hipDrvGraphAddMemsetNode.cc ) add_custom_target(add_Kernel.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/add_Kernel.cpp -o ${CMAKE_CURRENT_BINARY_DIR}/../graph/add_Kernel.code -I${HIP_PATH}/include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include --rocm-path=${ROCM_PATH}) diff --git a/projects/hip-tests/catch/unit/graph/graph_memset_node_test_common.hh b/projects/hip-tests/catch/unit/graph/graph_memset_node_test_common.hh index f4b957283e..b23a169339 100644 --- a/projects/hip-tests/catch/unit/graph/graph_memset_node_test_common.hh +++ b/projects/hip-tests/catch/unit/graph/graph_memset_node_test_common.hh @@ -26,14 +26,14 @@ THE SOFTWARE. #include #include -template void GraphMemsetNodeCommonPositive(F f) { +template void GraphMemsetNodeCommonPositive(F f) { const size_t width = GENERATE(1, 64, kPageSize / sizeof(T) + 1); const size_t height = GENERATE(1, 2, 1024); DYNAMIC_SECTION("Width: " << width << " Height: " << height) { LinearAllocGuard2D alloc(width, height); constexpr T set_value = 42; - hipMemsetParams params = {}; + Tp params = {}; params.dst = alloc.ptr(); params.elementSize = sizeof(T); params.width = width; @@ -50,7 +50,7 @@ template void GraphMemsetNodeCommonPositive(F f) { } } -template void MemsetCommonNegative(F f, hipMemsetParams params) { +template void MemsetCommonNegative(F f, T params) { SECTION("pMemsetParams == nullptr") { HIP_CHECK_ERROR(f(nullptr), hipErrorInvalidValue); } SECTION("pMemsetParams.dst == nullptr") { diff --git a/projects/hip-tests/catch/unit/graph/hipDrvGraphAddMemsetNode.cc b/projects/hip-tests/catch/unit/graph/hipDrvGraphAddMemsetNode.cc new file mode 100644 index 0000000000..a96988701d --- /dev/null +++ b/projects/hip-tests/catch/unit/graph/hipDrvGraphAddMemsetNode.cc @@ -0,0 +1,670 @@ +/* +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 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 + +#include +#include +#include +#include +#include + +#include "graph_memset_node_test_common.hh" +#include "graph_tests_common.hh" + + +/** + * @addtogroup hipDrvGraphAddMemsetNode hipDrvGraphAddMemsetNode + * @{ + * @ingroup GraphTest + * `hipDrvGraphAddMemsetNode(hipGraphNode_t* phGraphNode, hipGraph_t hGraph, const hipGraphNode_t* + * dependencies, size_t numDependencies, const HIP_MEMSET_NODE_PARAMS* memsetParams, hipCtx_t ctx)` + * - Creates a memset node and adds it to a graph + */ + +/** + * Test Description + * ------------------------ + * - Verify that all elements of destination memory are set to the correct value. + * The test is repeated for all valid element sizes(1, 2, 4), and several allocations of different + * height and width, both on host and device. + * Test source + * ------------------------ + * - unit/graph/hipDrvGraphAddMemsetNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEMPLATE_TEST_CASE("Unit_hipDrvGraphAddMemsetNode_Positive_Basic", "", uint8_t, uint16_t, + uint32_t) { + HIP_CHECK(hipInit(0)); + hipDevice_t device; + hipCtx_t context; + HIP_CHECK(hipDeviceGet(&device, 0)); + HIP_CHECK(hipCtxCreate(&context, 0, device)); + + CHECK_IMAGE_SUPPORT + + const auto f = [&context](HIP_MEMSET_NODE_PARAMS* params) { + hipGraph_t graph = nullptr; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + hipGraphNode_t node = nullptr; + HIP_CHECK(hipDrvGraphAddMemsetNode(&node, graph, nullptr, 0, params, context)); + + 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(hipGraphExecDestroy(graph_exec)); + HIP_CHECK(hipGraphDestroy(graph)); + + return hipSuccess; + }; + + GraphMemsetNodeCommonPositive(f); + + HIP_CHECK(hipCtxPopCurrent(&context)); + HIP_CHECK(hipCtxDestroy(context)); +} + +/** + * Test Description + * ------------------------ + * - Verify API behaviour with invalid arguments: + * -# pGraphNode 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 + * -# pMemsetParams is nullptr + * -# pMemsetParams::dst is nullptr + * -# pMemsetParams::elementSize is different from 1, 2, and 4 + * -# pMemsetParams::width is zero + * -# pMemsetParams::width is larger than the allocated memory region + * -# pMemsetParams::height is zero + * -# pMemsetParams::pitch is less than width when height is more than 1 + * -# pMemsetParams::pitch * pMemsetParams::height is larger than the allocated memory region + * Test source + * ------------------------ + * - unit/graph/hipDrvGraphAddMemsetNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipDrvGraphAddMemsetNode_Negative_Parameters") { + using namespace std::placeholders; + + 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)); + + LinearAllocGuard alloc(LinearAllocs::hipMalloc, 4 * sizeof(int)); + HIP_MEMSET_NODE_PARAMS params = {}; + params.dst = alloc.ptr(); + params.elementSize = sizeof(*alloc.ptr()); + params.width = 1; + params.height = 1; + params.value = 42; + + GraphAddNodeCommonNegativeTests( + std::bind(hipDrvGraphAddMemsetNode, _1, _2, _3, _4, ¶ms, context), graph); + + hipGraphNode_t node = nullptr; + MemsetCommonNegative(std::bind(hipDrvGraphAddMemsetNode, &node, graph, nullptr, 0, _1, context), + params); + + HIP_CHECK(hipGraphDestroy(graph)); + + HIP_CHECK(hipCtxPopCurrent(&context)); + HIP_CHECK(hipCtxDestroy(context)); +} + +/** + * Test Description + * ------------------------ + * - Allocate a 2D array using hipMallocPitch. Initialize the allocated memory using + * hipDrvGraphAddMemsetNode. Copy the values in device memory to host using + * hipDrvGraphAddMemcpyNode. Verify the results. + * Test source + * ------------------------ + * - unit/graph/hipDrvGraphAddMemsetNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipDrvGraphAddMemsetNode_hipMallocPitch_2D") { + HIP_CHECK(hipInit(0)); + hipDevice_t device; + hipCtx_t context; + HIP_CHECK(hipDeviceGet(&device, 0)); + HIP_CHECK(hipCtxCreate(&context, 0, device)); + + CHECK_IMAGE_SUPPORT + + size_t width = SIZE * sizeof(char), numW{SIZE}, numH{SIZE}, pitch_A; + char* A_d; + + hipGraph_t graph; + std::vector nodeDependencies; + // Host memory. + char* A_h = new char[numW * numH]; + for (size_t i = 0; i < numW; i++) { + for (size_t j = 0; j < numH; j++) { + *(A_h + i * numH + j) = ' '; + } + } + // 2D Memory allocation hipMallocPitch + HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), &pitch_A, width, numH)); + // Create Graph + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipGraphNode_t memsetNode, memcpyNode; + // Add MemSet Node + HIP_MEMSET_NODE_PARAMS memsetParams{}; + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(A_d); + memsetParams.value = memSetVal; + memsetParams.pitch = pitch_A; + memsetParams.elementSize = sizeof(char); + memsetParams.width = numW; + memsetParams.height = numH; + HIP_CHECK(hipDrvGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, &memsetParams, context)); + nodeDependencies.push_back(memsetNode); + + // Add MemCpy Node + auto srcPos = make_hipPos(0, 0, 0); + auto dstPos = make_hipPos(0, 0, 0); + auto srcPtr = make_hipPitchedPtr(A_d, pitch_A, numW, numH); + auto dstPtr = make_hipPitchedPtr(A_h, width, numW, numH); + auto extent = make_hipExtent(width, numH, 1); + hipMemcpyKind kind = hipMemcpyDeviceToHost; + + HIP_MEMCPY3D myparms = GetDrvMemcpy3DParms(dstPtr, dstPos, srcPtr, srcPos, extent, kind); + HIP_CHECK(hipDrvGraphAddMemcpyNode(&memcpyNode, graph, nodeDependencies.data(), + nodeDependencies.size(), &myparms, context)); + nodeDependencies.clear(); + // Create executable graph + hipStream_t streamForGraph; + hipGraphExec_t graphExec; + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Verfication + for (size_t i = 0; i < numW; i++) { + for (size_t j = 0; j < numH; j++) { + REQUIRE(*(A_h + i * numH + j) == memSetVal); + } + } + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + delete[] A_h; + HIP_CHECK(hipFree(A_d)); + + HIP_CHECK(hipCtxPopCurrent(&context)); + HIP_CHECK(hipCtxDestroy(context)); +} + +/** + * Test Description + * ------------------------ + * - Allocate a 1D array using hipMallocPitch. Initialize the allocated memory using + * hipDrvGraphAddMemsetNode. Copy the values in device memory to host using + * hipDrvGraphAddMemcpyNode. Verify the results. + * Test source + * ------------------------ + * - unit/graph/hipDrvGraphAddMemsetNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipDrvGraphAddMemsetNode_hipMallocPitch_1D") { + HIP_CHECK(hipInit(0)); + hipDevice_t device; + hipCtx_t context; + HIP_CHECK(hipDeviceGet(&device, 0)); + HIP_CHECK(hipCtxCreate(&context, 0, device)); + + CHECK_IMAGE_SUPPORT + + size_t width = SIZE * sizeof(char), numW{SIZE}, pitch_A; + char* A_d; + + // Initialize the host memory + std::vector A_h(numW, ' '); + + hipGraph_t graph; + std::vector nodeDependencies; + // 1D Memory allocation hipMallocPitch + HIP_CHECK(hipMallocPitch(reinterpret_cast(&A_d), &pitch_A, width, 1)); + // Create Graph + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipGraphNode_t memsetNode, memcpyNode; + // Add MemSet Node + HIP_MEMSET_NODE_PARAMS memsetParams{}; + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(A_d); + memsetParams.value = memSetVal; + memsetParams.pitch = pitch_A; + memsetParams.elementSize = sizeof(char); + memsetParams.width = numW; + memsetParams.height = 1; + HIP_CHECK(hipDrvGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, &memsetParams, context)); + nodeDependencies.push_back(memsetNode); + + // Add MemCpy Node + auto srcPos = make_hipPos(0, 0, 0); + auto dstPos = make_hipPos(0, 0, 0); + auto srcPtr = make_hipPitchedPtr(A_d, pitch_A, numW, 1); + auto dstPtr = make_hipPitchedPtr(A_h.data(), width, numW, 1); + auto extent = make_hipExtent(width, 1, 1); + hipMemcpyKind kind = hipMemcpyDeviceToHost; + + HIP_MEMCPY3D myparms = GetDrvMemcpy3DParms(dstPtr, dstPos, srcPtr, srcPos, extent, kind); + HIP_CHECK(hipDrvGraphAddMemcpyNode(&memcpyNode, graph, nodeDependencies.data(), + nodeDependencies.size(), &myparms, context)); + nodeDependencies.clear(); + + // Create executable graph + hipStream_t streamForGraph; + hipGraphExec_t graphExec; + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Verfication + for (size_t i = 0; i < numW; i++) { + REQUIRE(A_h[i] == memSetVal); + } + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + HIP_CHECK(hipFree(A_d)); + + HIP_CHECK(hipCtxPopCurrent(&context)); + HIP_CHECK(hipCtxDestroy(context)); +} + +/** + * Test Description + * ------------------------ + * - Allocate a 2D array using hipMalloc3D. Initialize the allocated memory using + * hipDrvGraphAddMemsetNode. Copy the values in device memory to host using + * hipDrvGraphAddMemcpyNode. Verify the results. + * Test source + * ------------------------ + * - unit/graph/hipDrvGraphAddMemsetNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipDrvGraphAddMemsetNode_hipMalloc3D_2D") { + HIP_CHECK(hipInit(0)); + hipDevice_t device; + hipCtx_t context; + HIP_CHECK(hipDeviceGet(&device, 0)); + HIP_CHECK(hipCtxCreate(&context, 0, device)); + + CHECK_IMAGE_SUPPORT + + size_t width = SIZE * sizeof(char); + size_t numW = SIZE, numH = SIZE; + + // Host Memory + char* A_h = new char[numW * numH]; + for (size_t i = 0; i < numW; i++) { + for (size_t j = 0; j < numH; j++) { + *(A_h + i * numH + j) = ' '; + } + } + hipGraph_t graph; + std::vector nodeDependencies; + + hipPitchedPtr A_d; + hipExtent extent3D = make_hipExtent(width, numH, 1); + + // Allocate 3D memory. + HIPCHECK(hipMalloc3D(&A_d, extent3D)); + + // Create Graph + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipGraphNode_t memsetNode, memcpyNode; + + // Add MemSet Node + HIP_MEMSET_NODE_PARAMS memsetParams{}; + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(A_d.ptr); + memsetParams.value = memSetVal; + memsetParams.pitch = A_d.pitch; + memsetParams.elementSize = sizeof(char); + memsetParams.width = numW; + memsetParams.height = numH; + HIP_CHECK(hipDrvGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, &memsetParams, context)); + nodeDependencies.push_back(memsetNode); + + // Add MemCpy Node + auto srcPos = make_hipPos(0, 0, 0); + auto dstPos = make_hipPos(0, 0, 0); + auto srcPtr = A_d; + auto dstPtr = make_hipPitchedPtr(A_h, width, numW, numH); + auto extent = make_hipExtent(width, numH, 1); + hipMemcpyKind kind = hipMemcpyDeviceToHost; + + HIP_MEMCPY3D myparms = GetDrvMemcpy3DParms(dstPtr, dstPos, srcPtr, srcPos, extent, kind); + HIP_CHECK(hipDrvGraphAddMemcpyNode(&memcpyNode, graph, nodeDependencies.data(), + nodeDependencies.size(), &myparms, context)); + nodeDependencies.clear(); + + // Create executable graph + hipStream_t streamForGraph; + hipGraphExec_t graphExec; + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Verfication + for (size_t i = 0; i < numW; i++) { + for (size_t j = 0; j < numH; j++) { + REQUIRE(*(A_h + i * numH + j) == memSetVal); + } + } + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + delete[] A_h; + HIP_CHECK(hipFree(A_d.ptr)); + + HIP_CHECK(hipCtxPopCurrent(&context)); + HIP_CHECK(hipCtxDestroy(context)); +} + +/** + * Test Description + * ------------------------ + * - Allocate a 1D array using hipMalloc3D. Initialize the allocated memory using + * hipDrvGraphAddMemsetNode. Copy the values in device memory to host using + * hipDrvGraphAddMemcpyNode. Verify the results. + * Test source + * ------------------------ + * - unit/graph/hipDrvGraphAddMemsetNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipDrvGraphAddMemsetNode_hipMalloc3D_1D") { + HIP_CHECK(hipInit(0)); + hipDevice_t device; + hipCtx_t context; + HIP_CHECK(hipDeviceGet(&device, 0)); + HIP_CHECK(hipCtxCreate(&context, 0, device)); + + CHECK_IMAGE_SUPPORT + + size_t width = SIZE * sizeof(char); + size_t numW = SIZE; + + // Initialize the host memory + std::vector A_h(numW, ' '); + + hipGraph_t graph; + std::vector nodeDependencies; + + hipPitchedPtr A_d; + hipExtent extent1D = make_hipExtent(width, 1, 1); + + // Allocate 3D memory. + HIPCHECK(hipMalloc3D(&A_d, extent1D)); + + // Create Graph + HIP_CHECK(hipGraphCreate(&graph, 0)); + hipGraphNode_t memsetNode, memcpyNode; + + // Add MemSet Node + HIP_MEMSET_NODE_PARAMS memsetParams{}; + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(A_d.ptr); + memsetParams.value = memSetVal; + memsetParams.pitch = A_d.pitch; + memsetParams.elementSize = sizeof(char); + memsetParams.width = numW; + memsetParams.height = 1; + HIP_CHECK(hipDrvGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, &memsetParams, context)); + nodeDependencies.push_back(memsetNode); + + // Add MemCpy Node + auto srcPos = make_hipPos(0, 0, 0); + auto dstPos = make_hipPos(0, 0, 0); + auto srcPtr = A_d; + auto dstPtr = make_hipPitchedPtr(A_h.data(), width, numW, 1); + auto extent = make_hipExtent(width, 1, 1); + hipMemcpyKind kind = hipMemcpyDeviceToHost; + + HIP_MEMCPY3D myparms = GetDrvMemcpy3DParms(dstPtr, dstPos, srcPtr, srcPos, extent, kind); + HIP_CHECK(hipDrvGraphAddMemcpyNode(&memcpyNode, graph, nodeDependencies.data(), + nodeDependencies.size(), &myparms, context)); + nodeDependencies.clear(); + + // Create executable graph + hipStream_t streamForGraph; + hipGraphExec_t graphExec; + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Verfication + for (size_t i = 0; i < numW; i++) { + REQUIRE(A_h[i] == memSetVal); + } + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)) + HIP_CHECK(hipFree(A_d.ptr)); + + HIP_CHECK(hipCtxPopCurrent(&context)); + HIP_CHECK(hipCtxDestroy(context)); +} + +/** + * Test Description + * ------------------------ + * - Allocate a 1D array using hipMalloc. Initialize the allocated memory using + * hipDrvGraphAddMemsetNode. Copy the values in device memory to host using + * hipDrvGraphAddMemcpyNode. Verify the results. + * Test source + * ------------------------ + * - unit/graph/hipDrvGraphAddMemsetNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipDrvGraphAddMemsetNode_hipMalloc_1D") { + HIP_CHECK(hipInit(0)); + hipDevice_t device; + hipCtx_t context; + HIP_CHECK(hipDeviceGet(&device, 0)); + HIP_CHECK(hipCtxCreate(&context, 0, device)); + + char* A_d; + size_t NumW = SIZE; + size_t Nbytes1D = SIZE * sizeof(char); + + // Initialize the host memory + std::vector A_h(NumW, ' '); + + // Allocate memory to Device pointer + HIP_CHECK(hipMalloc(reinterpret_cast(&A_d), Nbytes1D)); + + // Create the graph + hipGraph_t graph; + std::vector nodeDependencies; + hipGraphNode_t memsetNode, memcpyNode; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + // Add Memset node + HIP_MEMSET_NODE_PARAMS memsetParams{}; + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(A_d); + memsetParams.value = memSetVal; + memsetParams.pitch = Nbytes1D; + memsetParams.elementSize = sizeof(char); + memsetParams.width = NumW; + memsetParams.height = 1; + HIP_CHECK(hipDrvGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, &memsetParams, context)); + nodeDependencies.push_back(memsetNode); + + // Add MemCpy Node + hipPitchedPtr devPitchedPtr{A_d, Nbytes1D, NumW, 0}; + hipPitchedPtr hostPitchedPtr{A_h.data(), Nbytes1D, NumW, 0}; + auto srcPos = make_hipPos(0, 0, 0); + auto dstPos = make_hipPos(0, 0, 0); + auto srcPtr = devPitchedPtr; + auto dstPtr = hostPitchedPtr; + auto extent = make_hipExtent(Nbytes1D, 1, 1); + hipMemcpyKind kind = hipMemcpyDeviceToHost; + + HIP_MEMCPY3D myparms = GetDrvMemcpy3DParms(dstPtr, dstPos, srcPtr, srcPos, extent, kind); + HIP_CHECK(hipDrvGraphAddMemcpyNode(&memcpyNode, graph, nodeDependencies.data(), + nodeDependencies.size(), &myparms, context)); + nodeDependencies.clear(); + // Create executable graph + hipStream_t streamForGraph; + hipGraphExec_t graphExec; + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Verfication + for (size_t i = 0; i < NumW; i++) { + REQUIRE(A_h[i] == memSetVal); + } + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + HIP_CHECK(hipFree(A_d)); + + HIP_CHECK(hipCtxPopCurrent(&context)); + HIP_CHECK(hipCtxDestroy(context)); +} + +/** + * Test Description + * ------------------------ + * - Allocate a 1D array using hipMallocManaged. Initialize the allocated memory using + * hipDrvGraphAddMemsetNode. Copy the values in device memory to host using + * hipDrvGraphAddMemcpyNode. Verify the results. + * Test source + * ------------------------ + * - unit/graph/hipDrvGraphAddMemsetNode.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipGraphAddMemsetNode_hipMallocManaged") { + HIP_CHECK(hipInit(0)); + hipDevice_t device; + hipCtx_t context; + HIP_CHECK(hipDeviceGet(&device, 0)); + HIP_CHECK(hipCtxCreate(&context, 0, device)); + + int managed = 0; + HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory, 0)); + INFO("hipDeviceAttributeManagedMemory: " << managed); + if (managed != 1) { + WARN( + "GPU 0 doesn't support hipDeviceAttributeManagedMemory attribute" + "so defaulting to system memory."); + } + size_t Nbytes1D = SIZE * sizeof(char); + char* A_d; + // Initialize the host memory + std::vector A_h(SIZE, ' '); + // Device Memory + HIP_CHECK(hipMallocManaged(&A_d, SIZE * sizeof(char))); + // Create the graph + hipGraph_t graph; + std::vector nodeDependencies; + hipGraphNode_t memsetNode, memcpyNode; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + // Add Memset node + HIP_MEMSET_NODE_PARAMS memsetParams{}; + memset(&memsetParams, 0, sizeof(memsetParams)); + memsetParams.dst = reinterpret_cast(A_d); + memsetParams.value = memSetVal; + memsetParams.pitch = Nbytes1D; + memsetParams.elementSize = sizeof(char); + memsetParams.width = SIZE; + memsetParams.height = 1; + HIP_CHECK(hipDrvGraphAddMemsetNode(&memsetNode, graph, nullptr, 0, &memsetParams, context)); + nodeDependencies.push_back(memsetNode); + + // Add MemCpy Node + hipPitchedPtr devPitchedPtr{A_d, Nbytes1D, SIZE, 1}; + hipPitchedPtr hostPitchedPtr{A_h.data(), Nbytes1D, SIZE, 1}; + + auto srcPos = make_hipPos(0, 0, 0); + auto dstPos = make_hipPos(0, 0, 0); + auto srcPtr = devPitchedPtr; + auto dstPtr = hostPitchedPtr; + auto extent = make_hipExtent(Nbytes1D, 1, 1); + hipMemcpyKind kind = hipMemcpyDeviceToHost; + + HIP_MEMCPY3D myparms = GetDrvMemcpy3DParms(dstPtr, dstPos, srcPtr, srcPos, extent, kind); + HIP_CHECK(hipDrvGraphAddMemcpyNode(&memcpyNode, graph, nodeDependencies.data(), + nodeDependencies.size(), &myparms, context)); + nodeDependencies.clear(); + + // Create executable graph + hipStream_t streamForGraph; + hipGraphExec_t graphExec; + HIP_CHECK(hipStreamCreate(&streamForGraph)); + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, streamForGraph)); + HIP_CHECK(hipStreamSynchronize(streamForGraph)); + + // Verfication + for (size_t i = 0; i < SIZE; i++) { + REQUIRE(A_h[i] == memSetVal); + } + + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipStreamDestroy(streamForGraph)); + HIP_CHECK(hipFree(A_d)); + + HIP_CHECK(hipCtxPopCurrent(&context)); + HIP_CHECK(hipCtxDestroy(context)); +} diff --git a/projects/hip-tests/catch/unit/graph/hipGraphAddMemsetNode.cc b/projects/hip-tests/catch/unit/graph/hipGraphAddMemsetNode.cc index af502ab07a..e11d08a4b3 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphAddMemsetNode.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphAddMemsetNode.cc @@ -76,7 +76,7 @@ TEMPLATE_TEST_CASE("Unit_hipGraphAddMemsetNode_Positive_Basic", "", uint8_t, uin return hipSuccess; }; - GraphMemsetNodeCommonPositive(f); + GraphMemsetNodeCommonPositive(f); } /** diff --git a/projects/hip-tests/catch/unit/graph/hipGraphMemsetNodeSetParams.cc b/projects/hip-tests/catch/unit/graph/hipGraphMemsetNodeSetParams.cc index af8e6d50da..d816e8b88e 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphMemsetNodeSetParams.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphMemsetNodeSetParams.cc @@ -99,7 +99,7 @@ TEMPLATE_TEST_CASE("Unit_hipGraphMemsetNodeSetParams_Positive_Basic", "", uint8_ return hipSuccess; }; - GraphMemsetNodeCommonPositive(f); + GraphMemsetNodeCommonPositive(f); } /**