EXSWHTEC-374 - Implement tests for Graph device memory APIs #452

Change-Id: Ic82e3cda60296ec7b2acb9b5755c51ccbdb85c8d [ROCm/hip-tests commit: 27eeaa29a7]
2023-12-28 18:29:54 +01:00
commit 74299ca09d
@@ -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})
@@ -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 <hip_test_common.hh>
+#include <resource_guards.hh>
+#include <utils.hh>
+
+/**
+ * @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<int*>(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<void*>(&read_mem)));
+  REQUIRE(read_mem == used_mem);
+
+  attr = hipGraphMemAttrReservedMemCurrent;
+  HIP_CHECK(hipDeviceGetGraphMemAttribute(0, attr, reinterpret_cast<void*>(&read_mem)));
+  REQUIRE(read_mem == used_mem);
+
+  attr = hipGraphMemAttrUsedMemHigh;
+  HIP_CHECK(hipDeviceGetGraphMemAttribute(0, attr, reinterpret_cast<void*>(&read_mem)));
+  REQUIRE(read_mem == high_mem);
+
+  attr = hipGraphMemAttrReservedMemHigh;
+  HIP_CHECK(hipDeviceGetGraphMemAttribute(0, attr, reinterpret_cast<void*>(&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<void*>(&get_value)),
+        hipErrorInvalidDevice);
+  }
+
+  SECTION("Attribute value is not valid") {
+    HIP_CHECK_ERROR(hipDeviceGetGraphMemAttribute(0, static_cast<hipGraphMemAttributeType>(0x7),
+                                                  reinterpret_cast<void*>(&get_value)),
+                    hipErrorInvalidValue);
+  }
+
+  SECTION("Get value is nullptr") {
+    HIP_CHECK_ERROR(hipDeviceGetGraphMemAttribute(0, attr, nullptr), hipErrorInvalidValue);
+  }
+}
@@ -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 <hip_test_common.hh>
+
+/**
+ * @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);
+  }
+}
@@ -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 <hip_test_common.hh>
+
+/**
+ * @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<void*>(&set_value)),
+        hipErrorInvalidDevice);
+  }
+
+  SECTION("Attribute value is not supported") {
+    HIP_CHECK_ERROR(hipDeviceSetGraphMemAttribute(0, hipGraphMemAttrUsedMemCurrent,
+                                                  reinterpret_cast<void*>(&set_value)),
+                    hipErrorInvalidValue);
+  }
+
+  SECTION("Attribute value is not valid") {
+    HIP_CHECK_ERROR(hipDeviceSetGraphMemAttribute(0, static_cast<hipGraphMemAttributeType>(0x7),
+                                                  reinterpret_cast<void*>(&set_value)),
+                    hipErrorInvalidValue);
+  }
+
+  SECTION("Set hipGraphMemAttrUsedMemHigh to non-zero") {
+    size_t invalid_value = 1;
+    HIP_CHECK_ERROR(hipDeviceSetGraphMemAttribute(0, attr, reinterpret_cast<void*>(&invalid_value)),
+                    hipErrorInvalidValue);
+  }
+
+  SECTION("Set hipGraphMemAttrReservedMemHigh to non-zero") {
+    attr = hipGraphMemAttrReservedMemHigh;
+    size_t invalid_value = 1;
+    HIP_CHECK_ERROR(hipDeviceSetGraphMemAttribute(0, attr, reinterpret_cast<void*>(&invalid_value)),
+                    hipErrorInvalidValue);
+  }
+}