From dbc1a9d07383f3700e1da2b81144ef9e78adba94 Mon Sep 17 00:00:00 2001 From: Nives Vukovic Date: Sat, 9 Dec 2023 00:39:28 +0530 Subject: [PATCH] EXSWHTEC-36 - Implement tests for basic hipMemPool APIs (#437) Change-Id: Ice34feae7675699b04153d6fc3e10297428018b5 --- catch/include/hip_test_defgroups.hh | 7 ++ catch/unit/memory/CMakeLists.txt | 3 + catch/unit/memory/hipMemPoolCreate.cc | 93 ++++++++++++++ catch/unit/memory/hipMemPoolDestroy.cc | 71 +++++++++++ catch/unit/memory/hipMemPoolTrimTo.cc | 165 +++++++++++++++++++++++++ 5 files changed, 339 insertions(+) create mode 100644 catch/unit/memory/hipMemPoolCreate.cc create mode 100644 catch/unit/memory/hipMemPoolDestroy.cc create mode 100644 catch/unit/memory/hipMemPoolTrimTo.cc diff --git a/catch/include/hip_test_defgroups.hh b/catch/include/hip_test_defgroups.hh index c8636d4098..4f61381817 100644 --- a/catch/include/hip_test_defgroups.hh +++ b/catch/include/hip_test_defgroups.hh @@ -152,6 +152,13 @@ THE SOFTWARE. * @} */ +/** + * @defgroup StreamOTest Ordered Memory Allocator + * @{ + * This section describes the tests for Stream Ordered Memory Allocator functions of HIP runtime + * API. + */ + /** * @defgroup StreamTest Stream Management * @{ diff --git a/catch/unit/memory/CMakeLists.txt b/catch/unit/memory/CMakeLists.txt index 9432b2cb41..24e79fc6d9 100644 --- a/catch/unit/memory/CMakeLists.txt +++ b/catch/unit/memory/CMakeLists.txt @@ -109,6 +109,9 @@ set(TEST_SRC hipMemPoolApi.cc hipMemPoolSetGetAccess.cc hipMemPoolSetGetAttribute.cc + hipMemPoolCreate.cc + hipMemPoolDestroy.cc + hipMemPoolTrimTo.cc hipMemcpyPeer.cc hipMemcpyPeer_old.cc hipMemcpyPeerAsync.cc diff --git a/catch/unit/memory/hipMemPoolCreate.cc b/catch/unit/memory/hipMemPoolCreate.cc new file mode 100644 index 0000000000..57a3db7f45 --- /dev/null +++ b/catch/unit/memory/hipMemPoolCreate.cc @@ -0,0 +1,93 @@ +/* + 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 WARRANNTY OF ANY KIND, EXPRESS OR + IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER + LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + THE SOFTWARE. + */ + +#include + +/** + * @addtogroup hipMemPoolCreate hipMemPoolCreate + * @{ + * @ingroup StreamOTest + * `hipMemPoolCreate(hipMemPool_t* mem_pool, const hipMemPoolProps* pool_props)` - + * Creates a memory pool and returns the handle in mem pool + */ + +/** + * Test Description + * ------------------------ + * - Test to verify hipMemPoolCreate behavior with invalid arguments: + * -# Nullptr mem_pool + * -# Nullptr props + * -# Invalid props alloc type + * -# Invalid props location type + * -# Invalid props location id + * + * Test source + * ------------------------ + * - /unit/memory/hipMemPoolCreate.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipMemPoolCreate_Negative_Parameter") { + int mem_pool_support = 0; + HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, 0)); + if (!mem_pool_support) { + SUCCEED("Runtime doesn't support Memory Pool. Skip the test case."); + return; + } + + int num_dev = 0; + HIP_CHECK(hipGetDeviceCount(&num_dev)); + + hipMemPoolProps pool_props; + pool_props.allocType = hipMemAllocationTypePinned; + pool_props.handleTypes = hipMemHandleTypeNone; + pool_props.location.type = hipMemLocationTypeDevice; + pool_props.location.id = 0; + pool_props.win32SecurityAttributes = nullptr; + memset(pool_props.reserved, 0, sizeof(pool_props.reserved)); + + hipMemPool_t mem_pool = nullptr; + + SECTION("Passing nullptr to mem_pool") { + HIP_CHECK_ERROR(hipMemPoolCreate(nullptr, &pool_props), hipErrorInvalidValue); + } + + SECTION("Passing nullptr to props") { + HIP_CHECK_ERROR(hipMemPoolCreate(&mem_pool, nullptr), hipErrorInvalidValue); + } + + SECTION("Passing invalid props alloc type") { + pool_props.allocType = hipMemAllocationTypeInvalid; + HIP_CHECK_ERROR(hipMemPoolCreate(&mem_pool, &pool_props), hipErrorInvalidValue); + pool_props.allocType = hipMemAllocationTypePinned; + } + + SECTION("Passing invalid props location type") { + pool_props.location.type = hipMemLocationTypeInvalid; + HIP_CHECK_ERROR(hipMemPoolCreate(&mem_pool, &pool_props), hipErrorInvalidValue); + pool_props.location.type = hipMemLocationTypeDevice; + } + + SECTION("Passing invalid props location id") { + pool_props.location.id = num_dev; + HIP_CHECK_ERROR(hipMemPoolCreate(&mem_pool, &pool_props), hipErrorInvalidValue); + pool_props.location.id = 0; + } +} diff --git a/catch/unit/memory/hipMemPoolDestroy.cc b/catch/unit/memory/hipMemPoolDestroy.cc new file mode 100644 index 0000000000..f76059dbbc --- /dev/null +++ b/catch/unit/memory/hipMemPoolDestroy.cc @@ -0,0 +1,71 @@ +/* + 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 WARRANNTY OF ANY KIND, EXPRESS OR + IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER + LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + THE SOFTWARE. + */ + +#include "mempool_common.hh" + +/** + * @addtogroup hipMemPoolDestroy hipMemPoolDestroy + * @{ + * @ingroup StreamOTest + * `hipMemPoolDestroy(hipMemPool_t mem_pool)` - + * Destroys the specified memory pool + */ + +/** + * Test Description + * ------------------------ + * - Test to verify hipMemPoolCreate behavior with invalid arguments: + * -# Nullptr mem_pool + * -# Double hipMemPoolDestroy + * -# Attempt to destroy default mempool + * + * Test source + * ------------------------ + * - /unit/memory/hipMemPoolDestroy.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipMemPoolDestroy_Negative_Parameter") { + int mem_pool_support = 0; + HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, 0)); + if (!mem_pool_support) { + SUCCEED("Runtime doesn't support Memory Pool. Skip the test case."); + return; + } + + hipMemPool_t mem_pool = nullptr; + + SECTION("Passing nullptr to mempool") { + HIP_CHECK_ERROR(hipMemPoolDestroy(nullptr), hipErrorInvalidValue); + } + + SECTION("Double hipMemPoolDestroy") { + HIP_CHECK(hipMemPoolCreate(&mem_pool, &kPoolProps)); + HIP_CHECK(hipMemPoolDestroy(mem_pool)); + HIP_CHECK_ERROR(hipMemPoolDestroy(mem_pool), hipErrorInvalidValue); + } + + SECTION("Attempt to destroy default mempool") { + hipMemPool_t default_mem_pool = nullptr; + int device = 0; + HIP_CHECK(hipDeviceGetDefaultMemPool(&default_mem_pool, device)); + HIP_CHECK_ERROR(hipMemPoolDestroy(default_mem_pool), hipErrorInvalidValue); + } +} diff --git a/catch/unit/memory/hipMemPoolTrimTo.cc b/catch/unit/memory/hipMemPoolTrimTo.cc new file mode 100644 index 0000000000..6decc1c11f --- /dev/null +++ b/catch/unit/memory/hipMemPoolTrimTo.cc @@ -0,0 +1,165 @@ +/* + 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 WARRANNTY OF ANY KIND, EXPRESS OR + IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER + LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + THE SOFTWARE. + */ + + +#include "mempool_common.hh" + +#include +#include + +/** + * @addtogroup hipMemPoolTrimTo hipMemPoolTrimTo + * @{ + * @ingroup StreamOTest + * `hipMemPoolTrimTo(hipMemPool_t mem_pool, size_t min_bytes_to_hold)` - + * Releases freed memory back to the OS + */ + + +/** + * Test Description + * ------------------------ + * - Test to verify hipMemPoolTrimTo behavior with invalid arguments: + * -# Nullptr mem_pool + * + * Test source + * ------------------------ + * - /unit/memory/hipMemPoolTrimTo.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipMemPoolTrimTo_Negative_Parameter") { + int device_id = 0; + HIP_CHECK(hipSetDevice(device_id)); + + int mem_pool_support = 0; + HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, 0)); + if (!mem_pool_support) { + SUCCEED("Runtime doesn't support Memory Pool. Skip the test case."); + return; + } + + size_t trim_size = 1024; + + SECTION("Passing nullptr to mem_pool") { + HIP_CHECK_ERROR(hipMemPoolTrimTo(nullptr, trim_size), hipErrorInvalidValue); + } +} + + +/** + * Test Description + * ------------------------ + * - Basic test to verify hipMemPoolTrimTo releases memory correctly to the OS. + * Test source + * ------------------------ + * - /unit/memory/hipMemPoolTrimTo.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipMemPoolTrimTo_Positive_Basic") { + int device_id = 0; + HIP_CHECK(hipSetDevice(device_id)); + + int mem_pool_support = 0; + HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, 0)); + if (!mem_pool_support) { + SUCCEED("Runtime doesn't support Memory Pool. Skip the test case."); + return; + } + + const size_t allocation_size1 = kPageSize * kPageSize * 2; + const size_t allocation_size2 = kPageSize / 2; + MemPoolGuard mempool(MemPools::created, device_id); + + int* alloc_mem1; + int* alloc_mem2; + StreamGuard stream(Streams::created); + + HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast(&alloc_mem1), allocation_size1, + mempool.mempool(), stream.stream())); + HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast(&alloc_mem2), allocation_size2, + mempool.mempool(), stream.stream())); + + int blocks = 2; + int clk_rate; + if (IsGfx11()) { + HIP_CHECK(hipDeviceGetAttribute(&clk_rate, hipDeviceAttributeWallClockRate, 0)); + kernel_500ms_gfx11<<<32, blocks, 0, stream.stream()>>>(alloc_mem1, clk_rate); + } else { + HIP_CHECK(hipDeviceGetAttribute(&clk_rate, hipDeviceAttributeClockRate, 0)); + + kernel_500ms<<<32, blocks, 0, stream.stream()>>>(alloc_mem1, clk_rate); + } + + hipMemPoolAttr attr; + attr = hipMemPoolAttrReleaseThreshold; + // The pool must hold 128MB + std::uint64_t threshold = 128 * 1024 * 1024; + HIP_CHECK(hipMemPoolSetAttribute(mempool.mempool(), attr, &threshold)); + + // Not a real free, since kernel isn't done + HIP_CHECK(hipFreeAsync(reinterpret_cast(alloc_mem1), stream.stream())); + + // Get reserved memory before trim + attr = hipMemPoolAttrReservedMemCurrent; + std::uint64_t res_before_trim = 0; + HIP_CHECK(hipMemPoolGetAttribute(mempool.mempool(), attr, &res_before_trim)); + + size_t min_bytes_to_hold = allocation_size2; + HIP_CHECK(hipMemPoolTrimTo(mempool.mempool(), min_bytes_to_hold)); + + std::uint64_t res_after_trim = 0; + HIP_CHECK(hipMemPoolGetAttribute(mempool.mempool(), attr, &res_after_trim)); + // Trim must be a nop because execution isn't done + REQUIRE(res_before_trim == res_after_trim); + + HIP_CHECK(hipStreamSynchronize(stream.stream())); + + std::uint64_t res_after_sync = 0; + HIP_CHECK(hipMemPoolGetAttribute(mempool.mempool(), attr, &res_after_sync)); + // Since hipMemPoolAttrReleaseThreshold is 128 MB sync does nothing to the freed memory + REQUIRE(res_after_trim == res_after_sync); + + HIP_CHECK(hipMemPoolTrimTo(mempool.mempool(), min_bytes_to_hold)); + + HIP_CHECK(hipMemPoolGetAttribute(mempool.mempool(), attr, &res_after_trim)); + // Validate memory after real trim. The pool must hold less memory than before + REQUIRE(res_after_trim < res_after_sync); + + attr = hipMemPoolAttrReleaseThreshold; + std::uint64_t value64 = 0; + HIP_CHECK(hipMemPoolGetAttribute(mempool.mempool(), attr, &value64)); + // Make sure the threshold query works + REQUIRE(threshold == value64); + + attr = hipMemPoolAttrUsedMemCurrent; + HIP_CHECK(hipMemPoolGetAttribute(mempool.mempool(), attr, &value64)); + // Make sure the current usage query works - just small buffer left + REQUIRE(allocation_size2 == value64); + + attr = hipMemPoolAttrUsedMemHigh; + HIP_CHECK(hipMemPoolGetAttribute(mempool.mempool(), attr, &value64)); + // Make sure the high watermark usage works - the both buffers must be reported + REQUIRE((allocation_size1 + allocation_size2) == value64); + + HIP_CHECK(hipFreeAsync(reinterpret_cast(alloc_mem2), stream.stream())); +}