From aeae2fcbb1cc6779cd71d7593871b8c8e8aa9ced Mon Sep 17 00:00:00 2001 From: Jaydeep Patel Date: Mon, 22 Apr 2024 20:18:08 +0000 Subject: [PATCH] SWDEV-457316 - Fix tests to use graph's reserve memory instead of hipMemGetInfo. Change-Id: I839aaec3f502ed8686651fdc05b71c0da3e5dea2 [ROCm/hip-tests commit: b42adb24b47e715fc5022e3f73ccef521e32f870] --- .../catch/hipTestMain/config/config_amd_linux | 2 + .../hipTestMain/config/config_amd_windows | 2 + .../unit/graph/hipGraphAddMemAllocNode.cc | 799 +++++++++++++++++- .../unit/graph/hipGraphAddMemFreeNode.cc | 70 +- 4 files changed, 869 insertions(+), 4 deletions(-) diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux index ae37cd91f8..9ebbb88d5c 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux @@ -748,6 +748,8 @@ "Unit_hipGraphInstantiateWithFlags_AutoFreeOnLaunchDoubleKernel", "Unit_hipGraphInstantiateWithFlags_AutoFreeOnLaunchMultiProcess", "Unit_hipGraphInstantiateWithFlags_WithDefaultAndAutoFreeOnLaunch", + "=== SWDEV-457316 Below test is skipped due ref count logic (Discussed with German) ===", + "Unit_hipGraphAddMemAllocNode_Negative_Free_Alloc_Memory_Again", #endif #if defined gfx908 "=== Below test soft hang in stress test on 29/08/23 ===", diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows index db9c3ff75d..01c5cd65de 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows @@ -986,6 +986,8 @@ "Unit_hipGraphInstantiateWithFlags_AutoFreeOnLaunchFillKernel", "Unit_hipGraphInstantiateWithFlags_AutoFreeOnLaunchDoubleKernel", "Unit_hipGraphInstantiateWithFlags_WithDefaultAndAutoFreeOnLaunch", + "=== SWDEV-457316 Below test is skipped due ref count logic (Discussed with German) ===", + "Unit_hipGraphAddMemAllocNode_Negative_Free_Alloc_Memory_Again", #endif "=== Following tests disabled as it should be a local perf test", "Performance_hipExtLaunchKernelGGL_QueryGPUFrequency", diff --git a/projects/hip-tests/catch/unit/graph/hipGraphAddMemAllocNode.cc b/projects/hip-tests/catch/unit/graph/hipGraphAddMemAllocNode.cc index 64adc13600..346d1cab1a 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphAddMemAllocNode.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphAddMemAllocNode.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +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 @@ -21,6 +21,8 @@ THE SOFTWARE. */ #include +#include +#include #include #include @@ -278,7 +280,8 @@ static void createFreeGraph(hipGraphExec_t* graph_exec, int* device_alloc) { hipGraphNode_t free_node; HIP_CHECK(hipGraphCreate(&graph, 0)); - HIP_CHECK(hipGraphAddMemFreeNode(&free_node, graph, nullptr, 0, (void*)device_alloc)); + HIP_CHECK(hipGraphAddMemFreeNode(&free_node, graph, nullptr, 0, + (void*)device_alloc)); // Instantiate graph HIP_CHECK(hipGraphInstantiate(graph_exec, graph, nullptr, nullptr, 0)); @@ -457,6 +460,798 @@ TEST_CASE("Unit_hipGraphAddMemAllocNode_Positive_FreeSeparateGraph") { HIP_CHECK(hipGraphExecDestroy(graph_exec2)); } +/** +* Test Description +* ------------------------ +* - Create a graph and add 3 node with hipGraphAddMemAllocNode, +* initialize the memory allocated for 2 node. +* Add kernel node which will do vectorADD for these 2 allocated +* node and copy the result 3rd allocated node and verify. +* Test source +* ------------------------ +*  - /unit/graph/hipGraphAddMemAllocNode.cc +* Test requirements +* ------------------------ +*  - HIP_VERSION >= 6.1 +*/ + +TEST_CASE("Unit_hipGraphAddMemAllocNode_Functional_1") { + int mem_pool_support = 0; + HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, + hipDeviceAttributeMemoryPoolsSupported, 0)); + if (!mem_pool_support) { + HipTest::HIP_SKIP_TEST("Runtime doesn't support Memory Pool." + " Skip the test case."); + return; + } + + constexpr size_t N = 1024 * 1024; + constexpr size_t Nbytes = N * sizeof(int); + constexpr auto blocksPerCU = 6; // to hide latency + constexpr auto threadsPerBlock = 256; + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + size_t NElem{N}; + + hipGraph_t graph; + hipGraphExec_t graphExec; + hipStream_t stream; + hipGraphNode_t memcpyH2D_A, memcpyH2D_B, memcpyD2H_C; + hipGraphNode_t kernel_vecAdd; + hipKernelNodeParams kernelNodeParams{}; + hipGraphNode_t allocNodeA, freeNodeA, allocNodeB, freeNodeB; + hipGraphNode_t allocNodeC, freeNodeC; + hipMemAllocNodeParams allocParam; + + int *A_d, *B_d, *C_d; + int *A_h, *B_h, *C_h; + HipTest::initArrays(nullptr, nullptr, nullptr, + &A_h, &B_h, &C_h, N, false); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipStreamCreate(&stream)); + + memset(&allocParam, 0, sizeof(allocParam)); + allocParam.bytesize = Nbytes; + allocParam.poolProps.allocType = hipMemAllocationTypePinned; + allocParam.poolProps.location.id = 0; + allocParam.poolProps.location.type = hipMemLocationTypeDevice; + + HIP_CHECK(hipGraphAddMemAllocNode(&allocNodeA, graph, nullptr, + 0, &allocParam)); + REQUIRE(allocParam.dptr != nullptr); + A_d = reinterpret_cast(allocParam.dptr); + HIP_CHECK(hipGraphAddMemAllocNode(&allocNodeB, graph, + &allocNodeA, 1, &allocParam)); + REQUIRE(allocParam.dptr != nullptr); + B_d = reinterpret_cast(allocParam.dptr); + HIP_CHECK(hipGraphAddMemAllocNode(&allocNodeC, graph, + &allocNodeB, 1, &allocParam)); + REQUIRE(allocParam.dptr != nullptr); + C_d = reinterpret_cast(allocParam.dptr); + + // Check shows that A_d, B_d & C_d DON'T share any virtual address each other + REQUIRE(A_d != B_d); + REQUIRE(B_d != C_d); + REQUIRE(A_d != C_d); + + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, &allocNodeC, 1, A_d, + A_h, Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_B, graph, &allocNodeC, 1, B_d, + B_h, Nbytes, hipMemcpyHostToDevice)); + + void* kernelArgs[] = {&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(kernelArgs); + 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(hipGraphAddMemcpyNode1D(&memcpyD2H_C, graph, &kernel_vecAdd, 1, + C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); + + HIP_CHECK(hipGraphAddMemFreeNode(&freeNodeA, graph, &memcpyD2H_C, 1, + reinterpret_cast(A_d))); + HIP_CHECK(hipGraphAddMemFreeNode(&freeNodeB, graph, &memcpyD2H_C, 1, + reinterpret_cast(B_d))); + HIP_CHECK(hipGraphAddMemFreeNode(&freeNodeC, graph, &memcpyD2H_C, 1, + reinterpret_cast(C_d))); + + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + // Verify graph execution result + HipTest::checkVectorADD(A_h, B_h, C_h, N); + + HipTest::freeArrays(nullptr, nullptr, nullptr, A_h, B_h, C_h, false); + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipDeviceGraphMemTrim(0)); +} + +/** +* Test Description +* ------------------------ +* - Create a graph and add a node with hipGraphAddMemAllocNode and +* hipGraphAddMemFreeNode and launch it. Validate memory allocation + pointer shouldn't be null. +* Test source +* ------------------------ +*  - /unit/graph/hipGraphAddMemAllocNode.cc +* Test requirements +* ------------------------ +*  - HIP_VERSION >= 6.1 +*/ +TEST_CASE("Unit_hipGraphAddMemAllocNode_Functional_2") { + int mem_pool_support = 0; + HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, + hipDeviceAttributeMemoryPoolsSupported, 0)); + if (!mem_pool_support) { + HipTest::HIP_SKIP_TEST("Runtime doesn't support Memory Pool." + " Skip the test case."); + return; + } + + constexpr size_t Nbytes = 512 * 1024 *1024; + hipGraph_t graph; + hipGraphExec_t graphExec; + hipStream_t stream; + hipGraphNode_t allocNodeA, freeNodeA; + hipMemAllocNodeParams allocParam; + + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + + for ( int i = 0; i < numDevices; ++i ) { + HIP_CHECK(hipSetDevice(i)); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipStreamCreate(&stream)); + + memset(&allocParam, 0, sizeof(allocParam)); + allocParam.bytesize = Nbytes; + allocParam.poolProps.allocType = hipMemAllocationTypePinned; + allocParam.poolProps.location.id = 0; + allocParam.poolProps.location.type = hipMemLocationTypeDevice; + + HIP_CHECK(hipGraphAddMemAllocNode(&allocNodeA, graph, + NULL, 0, &allocParam)); + REQUIRE(allocParam.dptr != nullptr); + HIP_CHECK(hipGraphAddMemFreeNode(&freeNodeA, graph, &allocNodeA, 1, + reinterpret_cast(allocParam.dptr))); + + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + + size_t before = 0, after = 0; + HIP_CHECK(hipDeviceGraphMemTrim(i)); + HIP_CHECK(hipDeviceGetGraphMemAttribute(i, hipGraphMemAttrUsedMemCurrent, &before)); + + HIP_CHECK(hipGraphLaunch(graphExec, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + HIP_CHECK(hipDeviceGraphMemTrim(i)); + HIP_CHECK(hipDeviceGetGraphMemAttribute(i, hipGraphMemAttrUsedMemCurrent, &after)); + + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipStreamDestroy(stream)); + + REQUIRE(before == after); + } +} + +/** +* Test Description +* ------------------------ +* - Create a graph1 with hipGraphAddMemAllocNode and free it +* on different graph2 with hipGraphAddMemFreeNode. +* Test source +* ------------------------ +* - /unit/graph/hipGraphAddMemAllocNode.cc +* Test requirements +* ------------------------ +* - HIP_VERSION >= 6.1 +*/ +TEST_CASE("Unit_hipGraphAddMemAllocNode_Functional_3") { + int mem_pool_support = 0; + HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, + hipDeviceAttributeMemoryPoolsSupported, 0)); + if (!mem_pool_support) { + HipTest::HIP_SKIP_TEST("Runtime doesn't support Memory Pool." + " Skip the test case."); + return; + } + + constexpr size_t Nbytes = 512 * 1024 *1024; + hipGraph_t graph1, graph2; + hipGraphExec_t graphExec1, graphExec2; + hipStream_t stream; + hipGraphNode_t allocNodeA, freeNodeA; + hipMemAllocNodeParams allocParam; + + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + + for ( int i = 0; i < numDevices; i++ ) { + HIP_CHECK(hipSetDevice(i)); + + HIP_CHECK(hipGraphCreate(&graph1, 0)); + HIP_CHECK(hipGraphCreate(&graph2, 0)); + HIP_CHECK(hipStreamCreate(&stream)); + + memset(&allocParam, 0, sizeof(allocParam)); + allocParam.bytesize = Nbytes; + allocParam.poolProps.allocType = hipMemAllocationTypePinned; + allocParam.poolProps.location.id = 0; + allocParam.poolProps.location.type = hipMemLocationTypeDevice; + + HIP_CHECK(hipGraphAddMemAllocNode(&allocNodeA, graph1, + nullptr, 0, &allocParam)); + REQUIRE(allocParam.dptr != nullptr); + HIP_CHECK(hipGraphAddMemFreeNode(&freeNodeA, graph2, nullptr, 0, + reinterpret_cast(allocParam.dptr))); + + HIP_CHECK(hipGraphInstantiate(&graphExec1, graph1, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphInstantiate(&graphExec2, graph2, nullptr, nullptr, 0)); + + size_t before = 0, after = 0; + HIP_CHECK(hipDeviceGraphMemTrim(i)); + HIP_CHECK(hipDeviceGetGraphMemAttribute(i, hipGraphMemAttrUsedMemCurrent, &before)); + + HIP_CHECK(hipGraphLaunch(graphExec1, stream)); + HIP_CHECK(hipGraphLaunch(graphExec2, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + HIP_CHECK(hipDeviceGraphMemTrim(i)); + HIP_CHECK(hipDeviceGetGraphMemAttribute(i, hipGraphMemAttrUsedMemCurrent, &after)); + + HIP_CHECK(hipGraphDestroy(graph1)); + HIP_CHECK(hipGraphDestroy(graph2)); + HIP_CHECK(hipGraphExecDestroy(graphExec1)); + HIP_CHECK(hipGraphExecDestroy(graphExec2)); + HIP_CHECK(hipStreamDestroy(stream)); + + REQUIRE(before == after); + } +} + +/** +* Test Description +* ------------------------ +* - Create a graph1 with hipGraphAddMemAllocNode and free it with +* hipMemFreeAsync or hipMemFree. +* Test source +* ------------------------ +* - /unit/graph/hipGraphAddMemAllocNode.cc +* Test requirements +* ------------------------ +* - HIP_VERSION >= 6.1 +*/ +TEST_CASE("Unit_hipGraphAddMemAllocNode_Functional_4") { + int mem_pool_support = 0; + HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, + hipDeviceAttributeMemoryPoolsSupported, 0)); + if (!mem_pool_support) { + HipTest::HIP_SKIP_TEST("Runtime doesn't support Memory Pool." + " Skip the test case."); + return; + } + + constexpr size_t Nbytes = 512 * 1024 *1024; + hipGraph_t graph; + hipGraphExec_t graphExec; + hipStream_t stream; + hipGraphNode_t allocNodeA; + hipMemAllocNodeParams allocParam; + + void *temp; + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + + for ( int i = 0; i < numDevices; ++i ) { + HIP_CHECK(hipSetDevice(i)); + + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipStreamCreate(&stream)); + + memset(&allocParam, 0, sizeof(allocParam)); + allocParam.bytesize = Nbytes; + allocParam.poolProps.allocType = hipMemAllocationTypePinned; + allocParam.poolProps.location.id = 0; + allocParam.poolProps.location.type = hipMemLocationTypeDevice; + + HIP_CHECK(hipGraphAddMemAllocNode(&allocNodeA, graph, nullptr, + 0, &allocParam)); + temp = allocParam.dptr; + REQUIRE(temp != nullptr); + + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + + size_t before = 0, after = 0; + HIP_CHECK(hipDeviceGraphMemTrim(i)); + HIP_CHECK(hipDeviceGetGraphMemAttribute(i, hipGraphMemAttrUsedMemCurrent, &before)); + + HIP_CHECK(hipGraphLaunch(graphExec, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipDeviceGraphMemTrim(i)); + + HIP_CHECK(hipFree(temp)); + HIP_CHECK(hipDeviceGraphMemTrim(i)); + HIP_CHECK(hipDeviceGetGraphMemAttribute(i, hipGraphMemAttrUsedMemCurrent, &after)); + + REQUIRE(before == after); + } +} +/** +* Test Description +* ------------------------ +*  Negative Test for API hipGraphAddMemAllocNode - Argument validation check +1) Pass allocNode as nullptr. +2) Pass allocNode as empty structure. +3) Pass graph as nullptr. +4) Pass graph as empty structure. +5) Pass Dependencied node as null & number of dependencies as non zero. +6) Pass Dependencied node as valid & number of dependencies as 0 +7) Pass allocParam as nullptr. +8) Pass allocParam as empty structure. +9) Pass allocParam.bytesize as zero. +10) Pass allocParam.bytesize as INT_MAX. +11) Pass allocParam.poolProps.allocType as hipMemAllocationTypeInvalid. +12) Pass allocParam.poolProps.allocType as hipMemAllocationTypeMax. +13) Pass allocParam.poolProps.allocType as some unavailable id as -1 +14) Pass allocParam.poolProps.allocType as some unavailable id as INT_MAX +15) Pass allocParam.poolProps.location.id as -1. +16) Pass allocParam.poolProps.location.id as INT_MAX. +17) Pass allocParam.poolProps.location.type as hipMemLocationTypeInvalid +18) Pass allocParam.poolProps.location.type as hipMemLocationTypeDevice +19) Pass allocParam.poolProps.location.type as some unavailable id as -1 +20) Pass allocParam.poolProps.location.type as some unavailable id as INT_MAX +* Test source +* ------------------------ +*  - /unit/graph/hipGraphAddMemAllocNode.cc +* Test requirements +* ------------------------ +*  - HIP_VERSION >= 6.1 +*/ +TEST_CASE("Unit_hipGraphAddMemAllocNode_Argument_Check") { + int mem_pool_support = 0; + HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, + hipDeviceAttributeMemoryPoolsSupported, 0)); + if (!mem_pool_support) { + HipTest::HIP_SKIP_TEST("Runtime doesn't support Memory Pool." + " Skip the test case."); + return; + } + + constexpr size_t Nbytes = 512 * 1024 * 1024; + hipGraph_t graph; + hipGraphNode_t emptyNode, allocNodeA; + hipMemAllocNodeParams allocParam; + hipError_t ret; + + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipGraphAddEmptyNode(&emptyNode, graph, nullptr, 0)); + + memset(&allocParam, 0, sizeof(allocParam)); + allocParam.bytesize = Nbytes; + allocParam.poolProps.allocType = hipMemAllocationTypePinned; + allocParam.poolProps.location.id = 0; + allocParam.poolProps.location.type = hipMemLocationTypeDevice; + + HIP_CHECK(hipGraphAddMemAllocNode(&allocNodeA, graph, + nullptr, 0, &allocParam)); + REQUIRE(allocParam.dptr != nullptr); + + SECTION("Pass allocNode as nullptr.") { + ret = hipGraphAddMemAllocNode(nullptr, graph, nullptr, 0, &allocParam); + REQUIRE(ret == hipErrorInvalidValue); + } + SECTION("Pass allocNode as empty structure.") { + hipGraphNode_t allocNodeT{}; + ret = hipGraphAddMemAllocNode(&allocNodeT, graph, nullptr, + 0, &allocParam); + REQUIRE(ret == hipSuccess); + REQUIRE(allocParam.dptr != nullptr); + } + SECTION("Pass graph as nullptr.") { + ret = hipGraphAddMemAllocNode(&allocNodeA, nullptr, nullptr, + 0, &allocParam); + REQUIRE(ret == hipErrorInvalidValue); + } + SECTION("Pass graph as empty structure.") { + hipGraph_t graphT{}; + ret = hipGraphAddMemAllocNode(&allocNodeA, graphT, nullptr, + 0, &allocParam); + REQUIRE(ret == hipErrorInvalidValue); + } + SECTION("Dependencies node as nullptr & number of dependencies as nonzero") { + ret = hipGraphAddMemAllocNode(&allocNodeA, graph, nullptr, 1, &allocParam); + REQUIRE(ret == hipErrorInvalidValue); + } + SECTION("Dependencied node as valid & number of dependencies as 0") { + ret = hipGraphAddMemAllocNode(&allocNodeA, graph, &emptyNode, + 0, &allocParam); + REQUIRE(ret == hipSuccess); + REQUIRE(allocParam.dptr != nullptr); + } + SECTION("Pass allocParam as nullptr.") { + ret = hipGraphAddMemAllocNode(&allocNodeA, graph, nullptr, 0, nullptr); + REQUIRE(ret == hipErrorInvalidValue); + } +#if HT_NVIDIA + SECTION("Pass allocParam as empty structure.") { + hipMemAllocNodeParams allocParamT{}; + ret = hipGraphAddMemAllocNode(&allocNodeA, graph, nullptr, + 0, &allocParamT); + REQUIRE(ret == hipErrorInvalidValue); + } +#endif + SECTION("Pass allocParam.bytesize as INT_MAX.") { + memset(&allocParam, 0, sizeof(allocParam)); + allocParam.bytesize = INT_MAX; + allocParam.poolProps.allocType = hipMemAllocationTypePinned; + allocParam.poolProps.location.id = 0; + allocParam.poolProps.location.type = hipMemLocationTypeDevice; + + ret = hipGraphAddMemAllocNode(&allocNodeA, graph, nullptr, 0, &allocParam); + REQUIRE(ret == hipSuccess); + REQUIRE(allocParam.dptr != nullptr); + } +#if HT_NVIDIA + SECTION("Pass allocParam.bytesize as zero.") { + memset(&allocParam, 0, sizeof(allocParam)); + allocParam.bytesize = 0; + allocParam.poolProps.allocType = hipMemAllocationTypePinned; + allocParam.poolProps.location.id = 0; + allocParam.poolProps.location.type = hipMemLocationTypeDevice; + + ret = hipGraphAddMemAllocNode(&allocNodeA, graph, nullptr, 0, &allocParam); + REQUIRE(ret == hipErrorInvalidValue); + REQUIRE(allocParam.dptr == nullptr); + } + SECTION("Pass poolProps.allocType as hipMemAllocationTypeInvalid.") { + memset(&allocParam, 0, sizeof(allocParam)); + allocParam.bytesize = Nbytes; + allocParam.poolProps.allocType = hipMemAllocationTypeInvalid; + allocParam.poolProps.location.id = 0; + allocParam.poolProps.location.type = hipMemLocationTypeDevice; + + ret = hipGraphAddMemAllocNode(&allocNodeA, graph, nullptr, 0, &allocParam); + REQUIRE(allocParam.dptr == nullptr); + REQUIRE(ret == hipErrorInvalidValue); + } + SECTION("Pass allocParam.poolProps.allocType as hipMemAllocationTypeMax") { + memset(&allocParam, 0, sizeof(allocParam)); + allocParam.bytesize = Nbytes; + allocParam.poolProps.allocType = hipMemAllocationTypeMax; + allocParam.poolProps.location.id = 0; + allocParam.poolProps.location.type = hipMemLocationTypeDevice; + + ret = hipGraphAddMemAllocNode(&allocNodeA, graph, nullptr, 0, &allocParam); + REQUIRE(ret == hipErrorInvalidValue); + REQUIRE(allocParam.dptr == nullptr); + } + SECTION("Pass allocParam.poolProps.allocType id as -1") { + memset(&allocParam, 0, sizeof(allocParam)); + allocParam.bytesize = Nbytes; + allocParam.poolProps.allocType = hipMemAllocationType(-1); + allocParam.poolProps.location.id = 0; + allocParam.poolProps.location.type = hipMemLocationTypeDevice; + + ret = hipGraphAddMemAllocNode(&allocNodeA, graph, nullptr, 0, &allocParam); + REQUIRE(ret == hipErrorInvalidValue); + REQUIRE(allocParam.dptr == nullptr); + } + SECTION("Pass allocParam.poolProps.allocType id as INT_MAX") { + memset(&allocParam, 0, sizeof(allocParam)); + allocParam.bytesize = Nbytes; + allocParam.poolProps.allocType = hipMemAllocationType(INT_MAX); + allocParam.poolProps.location.id = 0; + allocParam.poolProps.location.type = hipMemLocationTypeDevice; + + ret = hipGraphAddMemAllocNode(&allocNodeA, graph, nullptr, 0, &allocParam); + REQUIRE(ret == hipErrorInvalidValue); + REQUIRE(allocParam.dptr == nullptr); + } + SECTION("Pass allocParam.poolProps.location.id as -1.") { + memset(&allocParam, 0, sizeof(allocParam)); + allocParam.bytesize = Nbytes; + allocParam.poolProps.allocType = hipMemAllocationTypePinned; + allocParam.poolProps.location.id = -1; + allocParam.poolProps.location.type = hipMemLocationTypeDevice; + + ret = hipGraphAddMemAllocNode(&allocNodeA, graph, nullptr, 0, &allocParam); + REQUIRE(ret == hipErrorInvalidValue); + REQUIRE(allocParam.dptr == nullptr); + } + SECTION("Pass allocParam.poolProps.location.id as INT_MAX.") { + memset(&allocParam, 0, sizeof(allocParam)); + allocParam.bytesize = Nbytes; + allocParam.poolProps.allocType = hipMemAllocationTypePinned; + allocParam.poolProps.location.id = INT_MAX; + allocParam.poolProps.location.type = hipMemLocationTypeDevice; + + ret = hipGraphAddMemAllocNode(&allocNodeA, graph, nullptr, 0, &allocParam); + REQUIRE(ret == hipErrorInvalidValue); + REQUIRE(allocParam.dptr == nullptr); + } + SECTION("allocParam.poolProps.location.type as hipMemLocationTypeInvalid") { + memset(&allocParam, 0, sizeof(allocParam)); + allocParam.bytesize = Nbytes; + allocParam.poolProps.allocType = hipMemAllocationTypePinned; + allocParam.poolProps.location.id = 0; + allocParam.poolProps.location.type = hipMemLocationTypeInvalid; + + ret = hipGraphAddMemAllocNode(&allocNodeA, graph, nullptr, 0, &allocParam); + REQUIRE(ret == hipErrorInvalidValue); + REQUIRE(allocParam.dptr == nullptr); + } +#endif + SECTION("allocParam.poolProps.location.type as hipMemLocationTypeDevice") { + memset(&allocParam, 0, sizeof(allocParam)); + allocParam.bytesize = Nbytes; + allocParam.poolProps.allocType = hipMemAllocationTypePinned; + allocParam.poolProps.location.id = 0; + allocParam.poolProps.location.type = hipMemLocationTypeDevice; + + ret = hipGraphAddMemAllocNode(&allocNodeA, graph, nullptr, 0, &allocParam); + REQUIRE(ret == hipSuccess); + REQUIRE(allocParam.dptr != nullptr); + } +#if HT_NVIDIA + SECTION("Pass allocParam.poolProps.location.type some unavailable id -1.") { + memset(&allocParam, 0, sizeof(allocParam)); + allocParam.bytesize = Nbytes; + allocParam.poolProps.allocType = hipMemAllocationTypePinned; + allocParam.poolProps.location.id = 0; + allocParam.poolProps.location.type = hipMemLocationType(-1); + + ret = hipGraphAddMemAllocNode(&allocNodeA, graph, nullptr, 0, &allocParam); + REQUIRE(ret == hipErrorInvalidValue); + REQUIRE(allocParam.dptr == nullptr); + } + SECTION("Pass allocParam.poolProps.location.type id INT_MAX") { + memset(&allocParam, 0, sizeof(allocParam)); + allocParam.bytesize = Nbytes; + allocParam.poolProps.allocType = hipMemAllocationTypePinned; + allocParam.poolProps.location.id = 0; + allocParam.poolProps.location.type = hipMemLocationType(INT_MAX); + + ret = hipGraphAddMemAllocNode(&allocNodeA, graph, nullptr, 0, &allocParam); + REQUIRE(ret == hipErrorInvalidValue); + REQUIRE(allocParam.dptr == nullptr); + } +#endif + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipDeviceGraphMemTrim(0)); +} + +/** +* Test Description +* ------------------------ +* - Negative Test for API hipGraphAddMemAllocNode & hipGraphAddMemFreeNode +* Create a graph with alloc and free node, Instanciate the graph twice and +* validate that instantiate should return error. +* Test source +* ------------------------ +* - /unit/graph/hipGraphAddMemAllocNode.cc +* Test requirements +* ------------------------ +* - HIP_VERSION >= 6.1 +*/ +TEST_CASE("Unit_hipGraphAddMemAllocNode_Negative_Instanciate_Graph_Again") { + int mem_pool_support = 0; + HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, + hipDeviceAttributeMemoryPoolsSupported, 0)); + if (!mem_pool_support) { + HipTest::HIP_SKIP_TEST("Runtime doesn't support Memory Pool." + " Skip the test case."); + return; + } + + constexpr size_t Nbytes = 512 * 1024 * 1024; + hipGraph_t graph; + hipGraphExec_t graphExec, graphExec1; + hipStream_t stream; + hipGraphNode_t allocNodeA, freeNodeA; + hipMemAllocNodeParams allocParam; + hipError_t ret; + + HIP_CHECK(hipDeviceGraphMemTrim(0)); + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipStreamCreate(&stream)); + + memset(&allocParam, 0, sizeof(allocParam)); + allocParam.bytesize = Nbytes; + allocParam.poolProps.allocType = hipMemAllocationTypePinned; + allocParam.poolProps.location.id = 0; + allocParam.poolProps.location.type = hipMemLocationTypeDevice; + + HIP_CHECK(hipGraphAddMemAllocNode(&allocNodeA, graph, + nullptr, 0, &allocParam)); + REQUIRE(allocParam.dptr != nullptr); + void *temp = allocParam.dptr; + + HIP_CHECK(hipGraphAddMemFreeNode(&freeNodeA, graph, &allocNodeA, 1, temp)); + + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + + // Instanciate 2nd time with same graph, it should give error + ret = hipGraphInstantiate(&graphExec1, graph, nullptr, nullptr, 0); + REQUIRE(ret == hipErrorNotSupported); + + HIP_CHECK(hipGraphLaunch(graphExec, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipDeviceGraphMemTrim(0)); +} + +/** +* Test Description +* ------------------------ +* - Negative Test for API hipGraphAddMemAllocNode & hipGraphAddMemFreeNode +* Create a graph with alloc and free node, try to free the same + alloc pointer again using hipFree(), it should return error. +* Test source +* ------------------------ +* - /unit/graph/hipGraphAddMemAllocNode.cc +* Test requirements +* ------------------------ +* - HIP_VERSION >= 6.1 +*/ +TEST_CASE("Unit_hipGraphAddMemAllocNode_Negative_Free_Alloc_Memory_Again") { + int mem_pool_support = 0; + HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, + hipDeviceAttributeMemoryPoolsSupported, 0)); + if (!mem_pool_support) { + HipTest::HIP_SKIP_TEST("Runtime doesn't support Memory Pool." + " Skip the test case."); + return; + } + + constexpr size_t Nbytes = 512 * 1024 * 1024; + hipGraph_t graph; + hipGraphExec_t graphExec; + hipStream_t stream; + hipGraphNode_t allocNodeA, freeNodeA; + hipMemAllocNodeParams allocParam; + hipError_t ret; + + HIP_CHECK(hipDeviceGraphMemTrim(0)); + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipStreamCreate(&stream)); + + memset(&allocParam, 0, sizeof(allocParam)); + allocParam.bytesize = Nbytes; + allocParam.poolProps.allocType = hipMemAllocationTypePinned; + allocParam.poolProps.location.id = 0; + allocParam.poolProps.location.type = hipMemLocationTypeDevice; + + HIP_CHECK(hipGraphAddMemAllocNode(&allocNodeA, graph, + nullptr, 0, &allocParam)); + REQUIRE(allocParam.dptr != nullptr); + void *temp = allocParam.dptr; + + HIP_CHECK(hipGraphAddMemFreeNode(&freeNodeA, graph, &allocNodeA, 1, temp)); + + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + + HIP_CHECK(hipGraphLaunch(graphExec, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + // Free alloc pointer manually again, it should give error + ret = hipFree(temp); + REQUIRE(ret == hipErrorInvalidValue); + + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipDeviceGraphMemTrim(0)); +} + +/** +* Test Description +* ------------------------ +*  Negative Test for API hipGraphAddMemAllocNode & hipGraphAddMemFreeNode + 1) Once free node added in any graph then we can't added it again. + 2) Clone graph should give error if a graph contain MemAlloc node. + 3) Clone graph should give error if a graph contain MemFree node. + 4) Destroy the MemAlloc node which was added in the graph. + 5) Destroy the MemFree node which was added in the graph. +* Test source +* ------------------------ +* - /unit/graph/hipGraphAddMemAllocNode.cc +* Test requirements +* ------------------------ +* - HIP_VERSION >= 6.1 +*/ +TEST_CASE("Unit_hipGraphAddMemAllocNode_Negative_With_Cloneed_Graph") { + int mem_pool_support = 0; + HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, + hipDeviceAttributeMemoryPoolsSupported, 0)); + if (!mem_pool_support) { + HipTest::HIP_SKIP_TEST("Runtime doesn't support Memory Pool." + " Skip the test case."); + return; + } + + constexpr size_t N = 512 * 1024 * 1024; + constexpr size_t Nbytes = N * sizeof(int); + hipGraph_t graph, graph_2, clone_graph, clone_graph_2; + hipGraphExec_t graphExec, graphExec_2; + hipStream_t stream; + hipGraphNode_t allocNodeA, freeNodeA, freeNodeB; + hipMemAllocNodeParams allocParam; + hipError_t ret; + + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipGraphCreate(&graph_2, 0)); + HIP_CHECK(hipStreamCreate(&stream)); + + memset(&allocParam, 0, sizeof(allocParam)); + allocParam.bytesize = Nbytes; + allocParam.poolProps.allocType = hipMemAllocationTypePinned; + allocParam.poolProps.location.id = 0; + allocParam.poolProps.location.type = hipMemLocationTypeDevice; + + HIP_CHECK(hipGraphAddMemAllocNode(&allocNodeA, graph, + nullptr, 0, &allocParam)); + REQUIRE(allocParam.dptr != nullptr); + void *temp = allocParam.dptr; + + HIP_CHECK(hipGraphAddMemFreeNode(&freeNodeA, graph_2, nullptr, 0, temp)); + + SECTION("Once free node added in any graph then we can't added it again") { + hipError_t ret1 = hipGraphAddMemFreeNode(&freeNodeB, graph, + nullptr, 0, temp); + REQUIRE(ret1 == hipErrorInvalidValue); + } + SECTION("Clone graph should give error if a graph contain memalloc node") { + ret = hipGraphClone(&clone_graph, graph); + REQUIRE(ret == hipErrorNotSupported); + } + SECTION("Clone graph should give error if a graph contain memfree node") { + ret = hipGraphClone(&clone_graph_2, graph_2); + REQUIRE(ret == hipErrorNotSupported); + } + SECTION("Destroy the MemAlloc node which was added in the graph") { + ret = hipGraphDestroyNode(allocNodeA); + REQUIRE(ret == hipErrorNotSupported); + } +#if HT_AMD + SECTION("Destroy the MemFree node which was added in the graph") { + ret = hipGraphDestroyNode(freeNodeB); + REQUIRE(ret == hipErrorInvalidValue); + } +#endif + + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec, stream)); + HIP_CHECK(hipGraphInstantiate(&graphExec_2, graph_2, nullptr, nullptr, 0)); + HIP_CHECK(hipGraphLaunch(graphExec_2, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipGraphDestroy(graph_2)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipGraphExecDestroy(graphExec_2)); + HIP_CHECK(hipStreamDestroy(stream)); +} /** * End doxygen group GraphTest. * @} diff --git a/projects/hip-tests/catch/unit/graph/hipGraphAddMemFreeNode.cc b/projects/hip-tests/catch/unit/graph/hipGraphAddMemFreeNode.cc index 42eafbc25e..518a248558 100644 --- a/projects/hip-tests/catch/unit/graph/hipGraphAddMemFreeNode.cc +++ b/projects/hip-tests/catch/unit/graph/hipGraphAddMemFreeNode.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +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 @@ -154,7 +154,8 @@ TEST_CASE("Unit_hipGraphAddMemFreeNode_Negative_NotSupported") { REQUIRE(alloc_param.dptr != nullptr); int* A_d = reinterpret_cast(alloc_param.dptr); - HIP_CHECK(hipGraphAddMemFreeNode(&free_node, graph2, nullptr, 0, (void*)A_d)); + HIP_CHECK(hipGraphAddMemFreeNode(&free_node, graph2, nullptr, 0, + (void*)A_d)); SECTION("More than one instantation of the graph exists") { hipGraphExec_t graph_exec1, graph_exec2; @@ -191,6 +192,71 @@ TEST_CASE("Unit_hipGraphAddMemFreeNode_Negative_NotSupported") { HIP_CHECK(hipGraphDestroy(graph2)); } + +/** +* Test Description +* ------------------------ +* - Functional Test for API hipGraphAddMemFreeNode - +* Measure memory footprint before creating graph. +* Create a graph and add a node with hipGraphAddMemAllocNode and +* hipGraphAddMemFreeNode and launch it. +* Measure memory footprint after the launch and destroy of the graph. +* Both before and after memory should be same after graph execution. +* Test source +* ------------------------ +* - /unit/graph/hipGraphAddMemFreeNode.cc +* Test requirements +* ------------------------ +* - HIP_VERSION >= 6.1 +*/ +TEST_CASE("Unit_hipGraphAddMemFreeNode_Functional") { + int mem_pool_support = 0; + HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, + hipDeviceAttributeMemoryPoolsSupported, 0)); + if (!mem_pool_support) { + HipTest::HIP_SKIP_TEST("Runtime doesn't support Memory Pool." + " Skip the test case."); + return; + } + + constexpr size_t Nbytes = 512 * 1024 *1024; + hipGraph_t graph; + hipGraphExec_t graphExec; + hipStream_t stream; + hipGraphNode_t allocNodeA, freeNodeA; + hipMemAllocNodeParams allocParam; + + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipStreamCreate(&stream)); + + memset(&allocParam, 0, sizeof(allocParam)); + allocParam.bytesize = Nbytes; + allocParam.poolProps.allocType = hipMemAllocationTypePinned; + allocParam.poolProps.location.id = 0; + allocParam.poolProps.location.type = hipMemLocationTypeDevice; + + HIP_CHECK(hipGraphAddMemAllocNode(&allocNodeA, graph, + NULL, 0, &allocParam)); + REQUIRE(allocParam.dptr != nullptr); + HIP_CHECK(hipGraphAddMemFreeNode(&freeNodeA, graph, &allocNodeA, 1, + reinterpret_cast(allocParam.dptr))); + + HIP_CHECK(hipGraphInstantiate(&graphExec, graph, nullptr, nullptr, 0)); + + size_t before = 0, after = 0; + HIP_CHECK(hipDeviceGraphMemTrim(0)); + HIP_CHECK(hipDeviceGetGraphMemAttribute(0, hipGraphMemAttrUsedMemCurrent, &before)); + HIP_CHECK(hipGraphLaunch(graphExec, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipDeviceGraphMemTrim(0)); + HIP_CHECK(hipDeviceGetGraphMemAttribute(0, hipGraphMemAttrUsedMemCurrent, &after)); + + HIP_CHECK(hipGraphDestroy(graph)); + HIP_CHECK(hipGraphExecDestroy(graphExec)); + HIP_CHECK(hipStreamDestroy(stream)); + + REQUIRE(before == after); +} /** * End doxygen group GraphTest. * @}