diff --git a/projects/hip-tests/catch/unit/graph/CMakeLists.txt b/projects/hip-tests/catch/unit/graph/CMakeLists.txt index c7b4036c49..d6519375a4 100644 --- a/projects/hip-tests/catch/unit/graph/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/graph/CMakeLists.txt @@ -154,6 +154,9 @@ set(TEST_SRC hipDrvGraphMemcpyNodeGetParams.cc hipDrvGraphMemcpyNodeSetParams.cc hipDrvGraphAddMemsetNode.cc + hipDeviceSetGraphMemAttribute.cc + hipDeviceGetGraphMemAttribute.cc + hipDeviceGraphMemTrim.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/hipDeviceGetGraphMemAttribute.cc b/projects/hip-tests/catch/unit/graph/hipDeviceGetGraphMemAttribute.cc new file mode 100644 index 0000000000..7aa10fe61c --- /dev/null +++ b/projects/hip-tests/catch/unit/graph/hipDeviceGetGraphMemAttribute.cc @@ -0,0 +1,205 @@ +/* +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 + +/** + * @addtogroup hipDeviceGetGraphMemAttribute hipDeviceGetGraphMemAttribute + * @{ + * @ingroup GraphTest + * `hipDeviceGetGraphMemAttribute(int device, hipGraphMemAttributeType attr, void* value)` - + * Get the mem attribute for graphs. + */ + +static constexpr auto element_count{64 * 1024 * 1024}; + + +/* Create graph with memory node */ +static void createGraph(hipGraphExec_t* graph_exec, int** device_alloc = nullptr) { + constexpr size_t num_bytes = element_count * sizeof(int); + + hipGraph_t graph; + HIP_CHECK(hipGraphCreate(&graph, 0)); + + hipGraphNode_t alloc_node; + hipMemAllocNodeParams alloc_param; + memset(&alloc_param, 0, sizeof(alloc_param)); + alloc_param.bytesize = num_bytes; + alloc_param.poolProps.allocType = hipMemAllocationTypePinned; + alloc_param.poolProps.location.id = 0; + alloc_param.poolProps.location.type = hipMemLocationTypeDevice; + + HIP_CHECK(hipGraphAddMemAllocNode(&alloc_node, graph, nullptr, 0, &alloc_param)); + REQUIRE(alloc_param.dptr != nullptr); + int* A_d = reinterpret_cast(alloc_param.dptr); + + if (device_alloc == nullptr) { + hipGraphNode_t free_node; + HIP_CHECK(hipGraphAddMemFreeNode(&free_node, graph, &alloc_node, 1, (void*)A_d)); + } else { + *device_alloc = A_d; + } + + // Instantiate graph + HIP_CHECK(hipGraphInstantiate(graph_exec, graph, nullptr, nullptr, 0)); + + HIP_CHECK(hipGraphDestroy(graph)); +} + +/* check if memory attributes for graphs contain expected values */ +static void checkGraphMemAttribute(size_t used_mem, size_t high_mem) { + size_t read_mem; + hipGraphMemAttributeType attr = hipGraphMemAttrUsedMemCurrent; + HIP_CHECK(hipDeviceGetGraphMemAttribute(0, attr, reinterpret_cast(&read_mem))); + REQUIRE(read_mem == used_mem); + + attr = hipGraphMemAttrReservedMemCurrent; + HIP_CHECK(hipDeviceGetGraphMemAttribute(0, attr, reinterpret_cast(&read_mem))); + REQUIRE(read_mem == used_mem); + + attr = hipGraphMemAttrUsedMemHigh; + HIP_CHECK(hipDeviceGetGraphMemAttribute(0, attr, reinterpret_cast(&read_mem))); + REQUIRE(read_mem == high_mem); + + attr = hipGraphMemAttrReservedMemHigh; + HIP_CHECK(hipDeviceGetGraphMemAttribute(0, attr, reinterpret_cast(&read_mem))); + REQUIRE(read_mem == high_mem); +} + +/** + * Test Description + * ------------------------ + * - Basic test to verify that hipDeviceGetGraphMemAttribute return correct memory attribute values + * when graphs with allocation nodes are launched, and after memory is freed to OS. + * Test source + * ------------------------ + * - /unit/graph/hipDeviceGetGraphMemAttribute.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipDeviceGetGraphMemAttribute_Positive_DoubleMemory") { + hipGraphExec_t graph_exec1, graph_exec2; + int *dev_p1, *dev_p2; + + StreamGuard stream_guard(Streams::created); + hipStream_t stream = stream_guard.stream(); + + createGraph(&graph_exec1, &dev_p1); + HIP_CHECK(hipGraphLaunch(graph_exec1, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + checkGraphMemAttribute(element_count * sizeof(int), element_count * sizeof(int)); + + createGraph(&graph_exec2, &dev_p2); + HIP_CHECK(hipGraphLaunch(graph_exec2, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + checkGraphMemAttribute(2 * element_count * sizeof(int), 2 * element_count * sizeof(int)); + + HIP_CHECK(hipFree(dev_p1)); + HIP_CHECK(hipFree(dev_p2)); + + HIP_CHECK(hipGraphExecDestroy(graph_exec1)); + HIP_CHECK(hipGraphExecDestroy(graph_exec2)); + HIP_CHECK(hipDeviceGraphMemTrim(0)); + checkGraphMemAttribute(0, 2 * element_count * sizeof(int)); +} + +/** + * Test Description + * ------------------------ + * - Basic test to verify that hipDeviceGetGraphMemAttribute return correct memory attribute values + * when graphs with allocation and free nodes are launched, and after memory is freed to OS. + * Test source + * ------------------------ + * - /unit/graph/hipDeviceGetGraphMemAttribute.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipDeviceGetGraphMemAttribute_Positive_ReuseMemory") { + hipGraphExec_t graph_exec1, graph_exec2; + + StreamGuard stream_guard(Streams::created); + hipStream_t stream = stream_guard.stream(); + + createGraph(&graph_exec1); + HIP_CHECK(hipGraphLaunch(graph_exec1, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + checkGraphMemAttribute(element_count * sizeof(int), element_count * sizeof(int)); + + createGraph(&graph_exec2); + HIP_CHECK(hipGraphLaunch(graph_exec2, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + + checkGraphMemAttribute(element_count * sizeof(int), element_count * sizeof(int)); + + HIP_CHECK(hipGraphExecDestroy(graph_exec1)); + HIP_CHECK(hipGraphExecDestroy(graph_exec2)); + HIP_CHECK(hipDeviceGraphMemTrim(0)); + checkGraphMemAttribute(0, element_count * sizeof(int)); +} + +/** + * Test Description + * ------------------------ + * - Test to verify hipDeviceGetGraphMemAttribute behavior with invalid arguments: + * -# Device is not valid + * -# Attribute value is not valid + * -# Get value is nullptr + * Test source + * ------------------------ + * - /unit/graph/hipDeviceGetGraphMemAttribute.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipDeviceGetGraphMemAttribute_Negative_Parameters") { + int device_id = 0; + HIP_CHECK(hipSetDevice(device_id)); + + int num_dev = 0; + HIP_CHECK(hipGetDeviceCount(&num_dev)); + + hipGraphMemAttributeType attr = hipGraphMemAttrUsedMemHigh; + size_t get_value = 0; + + SECTION("Device is not valid") { + HIP_CHECK_ERROR( + hipDeviceGetGraphMemAttribute(num_dev, attr, reinterpret_cast(&get_value)), + hipErrorInvalidDevice); + } + + SECTION("Attribute value is not valid") { + HIP_CHECK_ERROR(hipDeviceGetGraphMemAttribute(0, static_cast(0x7), + reinterpret_cast(&get_value)), + hipErrorInvalidValue); + } + + SECTION("Get value is nullptr") { + HIP_CHECK_ERROR(hipDeviceGetGraphMemAttribute(0, attr, nullptr), hipErrorInvalidValue); + } +} diff --git a/projects/hip-tests/catch/unit/graph/hipDeviceGraphMemTrim.cc b/projects/hip-tests/catch/unit/graph/hipDeviceGraphMemTrim.cc new file mode 100644 index 0000000000..5d730a3f88 --- /dev/null +++ b/projects/hip-tests/catch/unit/graph/hipDeviceGraphMemTrim.cc @@ -0,0 +1,73 @@ +/* +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 + +/** + * @addtogroup hipDeviceGraphMemTrim hipDeviceGraphMemTrim + * @{ + * @ingroup GraphTest + * `hipDeviceGraphMemTrim(int device)` - Free unused memory on specific device used for graph back + * to OS. + */ + +/** + * Test Description + * ------------------------ + * - Basic test to verify that unused memory used for graph can be freed on each device. + * Test source + * ------------------------ + * - /unit/graph/hipDeviceGraphMemTrim.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipDeviceGraphMemTrim_Positive_Default") { + const auto device = GENERATE(range(0, HipTest::getDeviceCount())); + + // Check for each device + HIP_CHECK(hipDeviceGraphMemTrim(device)); +} + +/** + * Test Description + * ------------------------ + * - Test to verify hipDeviceGraphMemTrim behavior with invalid arguments: + * -# Device is not valid + * Test source + * ------------------------ + * - /unit/graph/hipDeviceGraphMemTrim.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipDeviceGraphMemTrim_Negative_Parameters") { + int device_id = 0; + HIP_CHECK(hipSetDevice(device_id)); + + int num_dev = 0; + HIP_CHECK(hipGetDeviceCount(&num_dev)); + + SECTION("Device is not valid") { + HIP_CHECK_ERROR(hipDeviceGraphMemTrim(num_dev), hipErrorInvalidDevice); + } +} diff --git a/projects/hip-tests/catch/unit/graph/hipDeviceSetGraphMemAttribute.cc b/projects/hip-tests/catch/unit/graph/hipDeviceSetGraphMemAttribute.cc new file mode 100644 index 0000000000..a103b12fee --- /dev/null +++ b/projects/hip-tests/catch/unit/graph/hipDeviceSetGraphMemAttribute.cc @@ -0,0 +1,117 @@ +/* +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 + +/** + * @addtogroup hipDeviceSetGraphMemAttribute hipDeviceSetGraphMemAttribute + * @{ + * @ingroup GraphTest + * `hipDeviceSetGraphMemAttribute(int device, hipGraphMemAttributeType attr, void* value)` - + * Set the mem attribute for graphs. + */ + +static void GraphSetGetAttribute(int device, hipGraphMemAttributeType attr, size_t set_value) { + size_t get_value = 100; + HIP_CHECK(hipDeviceSetGraphMemAttribute(device, attr, &set_value)); + HIP_CHECK(hipDeviceGetGraphMemAttribute(device, attr, &get_value)); + REQUIRE(get_value == set_value); +} + +/** + * Test Description + * ------------------------ + * - Basic test to verify that valid attributes can be reset to zero. + * Test source + * ------------------------ + * - /unit/graph/hipDeviceSetGraphMemAttribute.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipDeviceSetGraphMemAttribute_Positive_Default") { + const auto device = GENERATE(range(0, HipTest::getDeviceCount())); + const auto attr_type = GENERATE(hipGraphMemAttrUsedMemHigh, hipGraphMemAttrReservedMemHigh); + + // Check if attributes can be reset + size_t set_value = 0; + GraphSetGetAttribute(device, attr_type, set_value); +} + + +/** + * Test Description + * ------------------------ + * - Test to verify hipDeviceSetGraphMemAttribute behavior with invalid arguments: + * -# Device is not valid + * -# Attribute value is not supported + * -# Attribute value is not valid + * -# Set hipGraphMemAttrUsedMemHigh to non-zero + * -# Set hipGraphMemAttrReservedMemHigh to non-zero + * Test source + * ------------------------ + * - /unit/graph/hipDeviceSetGraphMemAttribute.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipDeviceSetGraphMemAttribute_Negative_Parameters") { + int device_id = 0; + HIP_CHECK(hipSetDevice(device_id)); + + int num_dev = 0; + HIP_CHECK(hipGetDeviceCount(&num_dev)); + + hipGraphMemAttributeType attr = hipGraphMemAttrUsedMemHigh; + size_t set_value = 0; + + SECTION("device is not valid") { + HIP_CHECK_ERROR( + hipDeviceSetGraphMemAttribute(num_dev, attr, reinterpret_cast(&set_value)), + hipErrorInvalidDevice); + } + + SECTION("Attribute value is not supported") { + HIP_CHECK_ERROR(hipDeviceSetGraphMemAttribute(0, hipGraphMemAttrUsedMemCurrent, + reinterpret_cast(&set_value)), + hipErrorInvalidValue); + } + + SECTION("Attribute value is not valid") { + HIP_CHECK_ERROR(hipDeviceSetGraphMemAttribute(0, static_cast(0x7), + reinterpret_cast(&set_value)), + hipErrorInvalidValue); + } + + SECTION("Set hipGraphMemAttrUsedMemHigh to non-zero") { + size_t invalid_value = 1; + HIP_CHECK_ERROR(hipDeviceSetGraphMemAttribute(0, attr, reinterpret_cast(&invalid_value)), + hipErrorInvalidValue); + } + + SECTION("Set hipGraphMemAttrReservedMemHigh to non-zero") { + attr = hipGraphMemAttrReservedMemHigh; + size_t invalid_value = 1; + HIP_CHECK_ERROR(hipDeviceSetGraphMemAttribute(0, attr, reinterpret_cast(&invalid_value)), + hipErrorInvalidValue); + } +}