diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux index 2b93de19f8..983bf2b0a8 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux @@ -127,7 +127,6 @@ "Unit_deviceAllocation_InOneThread_AccessInAllThreads", "=== Patch which removes the typetraits implementation from std namespace in hiprtc is reverted ===", "Unit_hiprtc_stdheaders", -<<<<<<< HEAD "Unit_hipMemAddressFree_negative", "Unit_hipMemAddressReserve_AlignmentTest", "Unit_hipMemAddressReserve_Negative", @@ -250,7 +249,6 @@ "Unit_hipGraphicsResourceGetMappedPointer_Negative_Parameters", "Unit_hipGraphicsUnmapResources_Negative_Parameters", "Unit_hipGraphicsUnregisterResource_Negative_Parameters", -======= "=== Below tests fail in external CI for PR https://github.com/ROCm-Developer-Tools/hip-tests/pull/356 ===", "Unit_Device_Complex_Unary_Negative_Parameters_RTC", "Unit_Device_Complex_Binary_Negative_Parameters_RTC", @@ -264,7 +262,6 @@ "Unit_Device_Complex_Binary_float_Negative", "Unit_Device_Complex_Binary_double_Negative", "Unit_Device_Complex_hipCfma_Negative", ->>>>>>> ce6de407 (EXSWHTEC-316 - Implement tests for Complex type functions (#356)) #endif #if defined VEGA20 "=== SWDEV-419112 Below tests fail in stress test on 29/08/23 ===", diff --git a/projects/hip-tests/catch/hipTestMain/config/config_nvidia_linux.json b/projects/hip-tests/catch/hipTestMain/config/config_nvidia_linux.json index d2467d0fca..ccf816ea0f 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_nvidia_linux.json +++ b/projects/hip-tests/catch/hipTestMain/config/config_nvidia_linux.json @@ -19,7 +19,7 @@ "Unit_ChannelDescriptor_Positive_Basic_4D - ulong4", "Unit_ChannelDescriptor_Positive_Basic_4D - long4", "=== Below test fails in external CI for PR https://github.com/ROCm-Developer-Tools/hip-tests/pull/38 ===", - "Unit_hipFreeAsync_negative", + "Unit_hipFreeAsync_Negative_Parameters", "=== Below test fails in external CI for PR https://github.com/ROCm-Developer-Tools/hip-tests/pull/92 ===", "Unit_hipGetTexObjectResourceDesc_positive", "Unit_hipGetTexObjectResourceDesc_Negative_Parameters", diff --git a/projects/hip-tests/catch/include/hip_test_defgroups.hh b/projects/hip-tests/catch/include/hip_test_defgroups.hh index cb95778f3e..fdf289e8c5 100644 --- a/projects/hip-tests/catch/include/hip_test_defgroups.hh +++ b/projects/hip-tests/catch/include/hip_test_defgroups.hh @@ -159,6 +159,13 @@ THE SOFTWARE. * API. */ +/** + * @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/projects/hip-tests/catch/unit/memory/CMakeLists.txt b/projects/hip-tests/catch/unit/memory/CMakeLists.txt index 90d0690bf0..09e7277aac 100644 --- a/projects/hip-tests/catch/unit/memory/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/memory/CMakeLists.txt @@ -112,6 +112,7 @@ set(TEST_SRC hipMemPoolCreate.cc hipMemPoolDestroy.cc hipMemPoolTrimTo.cc + hipMallocFromPoolAsync.cc hipMemcpyPeer.cc hipMemcpyPeer_old.cc hipMemcpyPeerAsync.cc diff --git a/projects/hip-tests/catch/unit/memory/hipFreeAsync.cc b/projects/hip-tests/catch/unit/memory/hipFreeAsync.cc index c7fcd0b179..eb8c5b50de 100644 --- a/projects/hip-tests/catch/unit/memory/hipFreeAsync.cc +++ b/projects/hip-tests/catch/unit/memory/hipFreeAsync.cc @@ -21,24 +21,59 @@ THE SOFTWARE. */ #include +#include +#include -TEST_CASE("Unit_hipFreeAsync_negative") { - HIP_CHECK(hipSetDevice(0)); - void* p = nullptr; - hipStream_t stream{nullptr}; - HIP_CHECK(hipStreamCreate(&stream)); +/** + * @addtogroup hipFreeAsync hipFreeAsync + * @{ + * @ingroup StreamOTest + * `hipFreeAsync(void* dev_ptr, hipStream_t stream)` + * - Frees memory with stream ordered semantics + */ - SECTION("dev_ptr is nullptr") { REQUIRE(hipFreeAsync(nullptr, stream) != hipSuccess); } - SECTION("invalid stream handle") { - HIP_CHECK(hipMallocAsync(static_cast(&p), 100, stream)); - HIP_CHECK(hipStreamSynchronize(stream)); - hipError_t error = hipFreeAsync(p, reinterpret_cast(-1)); - HIP_CHECK(hipFreeAsync(p, stream)); - HIP_CHECK(hipStreamSynchronize(stream)); - REQUIRE(error != hipSuccess); +/** + * Test Description + * ------------------------ + * - Test to verify hipFreeAsync behavior with invalid arguments: + * -# Nullptr dev_ptr + * -# Invalid stream handle + * -# Double hipFreeAsync + * + * Test source + * ------------------------ + * - /unit/memory/hipFreeAsync.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipFreeAsync_Negative_Parameters") { + int device_id = 0; + HIP_CHECK(hipSetDevice(device_id)); + + int* p = nullptr; + size_t alloc_size = 1024; + StreamGuard stream(Streams::created); + + SECTION("dev_ptr is nullptr") { + HIP_CHECK_ERROR(hipFreeAsync(nullptr, stream.stream()), hipErrorInvalidValue); } - HIP_CHECK(hipStreamSynchronize(stream)); - HIP_CHECK(hipStreamDestroy(stream)); + SECTION("Invalid stream handle") { + HIP_CHECK(hipMallocAsync(reinterpret_cast(&p), alloc_size, stream.stream())); + HIP_CHECK(hipStreamSynchronize(stream.stream())); + HIP_CHECK_ERROR(hipFreeAsync(p, reinterpret_cast(-1)), hipErrorInvalidHandle); + HIP_CHECK(hipFreeAsync(reinterpret_cast(p), stream.stream())); + HIP_CHECK(hipStreamSynchronize(stream.stream())); + } + + SECTION("Double free") { + HIP_CHECK(hipMallocAsync(reinterpret_cast(&p), alloc_size, stream.stream())); + HIP_CHECK(hipStreamSynchronize(stream.stream())); + HIP_CHECK(hipFreeAsync(reinterpret_cast(p), stream.stream())); + HIP_CHECK(hipStreamSynchronize(stream.stream())); + HIP_CHECK_ERROR(hipFreeAsync(reinterpret_cast(p), stream.stream()), + hipErrorInvalidValue); + } } diff --git a/projects/hip-tests/catch/unit/memory/hipMallocAsync.cc b/projects/hip-tests/catch/unit/memory/hipMallocAsync.cc index db74c99459..f9e1d8aa8a 100644 --- a/projects/hip-tests/catch/unit/memory/hipMallocAsync.cc +++ b/projects/hip-tests/catch/unit/memory/hipMallocAsync.cc @@ -17,31 +17,122 @@ THE SOFTWARE. */ -#include -#include +#include "mempool_common.hh" #include -TEST_CASE("Unit_hipMallocAsync_negative") { - HIP_CHECK(hipSetDevice(0)); +#pragma clang diagnostic ignored "-Wunused-parameter" - void* p = nullptr; - size_t max_size = std::numeric_limits::max(); - hipStream_t stream{nullptr}; - HIP_CHECK(hipStreamCreate(&stream)); +/** + * @addtogroup hipMallocAsync hipMallocAsync + * @{ + * @ingroup StreamOTest + * `hipMallocAsync(void** dev_ptr, size_t size, hipStream_t stream)` + * - Allocates memory with stream ordered semantics + */ - - SECTION("Device pointer is null") { REQUIRE(hipMallocAsync(nullptr, 100, stream) != hipSuccess); } - - SECTION("stream is invalid") { - REQUIRE(hipMallocAsync(static_cast(&p), 100, reinterpret_cast(-1)) != - hipSuccess); - } - - SECTION("out of memory") { - REQUIRE(hipMallocAsync(static_cast(&p), max_size, stream) != hipSuccess); - } - - HIP_CHECK(hipStreamSynchronize(stream)); - HIP_CHECK(hipStreamDestroy(stream)); +/** + * Test Description + * ------------------------ + * - Basic test to verify proper allocation and stream ordering of hipMallocAsync when one + * memory allocation is performed. + * Test source + * ------------------------ + * - /unit/memory/hipMallocAsync.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipMallocAsync_Basic_OneAlloc") { + MallocMemPoolAsync_OneAlloc( + [](void** dev_ptr, size_t size, hipMemPool_t mem_pool, hipStream_t stream) { + return hipMallocAsync(dev_ptr, size, stream); + }, + MemPools::dev_default); +} + +/** + * Test Description + * ------------------------ + * - Basic test to verify proper allocation and stream ordering of hipMallocAsync when two + * memory allocations are performed. + * Test source + * ------------------------ + * - /unit/memory/hipMallocAsync.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipMallocAsync_Basic_TwoAllocs") { + MallocMemPoolAsync_TwoAllocs( + [](void** dev_ptr, size_t size, hipMemPool_t mem_pool, hipStream_t stream) { + return hipMallocAsync(dev_ptr, size, stream); + }, + MemPools::dev_default); +} + +/** + * Test Description + * ------------------------ + * - Basic test to verify that memory allocated with hipMallocAsync can be properly reused. + * Test source + * ------------------------ + * - /unit/memory/hipMallocAsync.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipMallocAsync_Basic_Reuse") { + MallocMemPoolAsync_Reuse([](void** dev_ptr, size_t size, hipMemPool_t mem_pool, + hipStream_t stream) { return hipMallocAsync(dev_ptr, size, stream); }, + MemPools::dev_default); +} + + +/** + * Test Description + * ------------------------ + * - Test to verify hipMallocAsync behavior with invalid arguments: + * -# Nullptr dev_ptr + * -# Invalid stream handle + * -# Size is max size_t + * + * Test source + * ------------------------ + * - /unit/memory/hipMallocAsync.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipMallocAsync_Negative_Parameters") { + 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; + } + + int* p = nullptr; + size_t max_size = std::numeric_limits::max(); + size_t alloc_size = 1024; + MemPoolGuard mempool(MemPools::dev_default, device_id); + StreamGuard stream(Streams::created); + + SECTION("dev_ptr is nullptr") { + HIP_CHECK_ERROR(hipMallocAsync(nullptr, alloc_size, stream.stream()), hipErrorInvalidValue); + } + + SECTION("invalid stream handle") { + HIP_CHECK_ERROR( + hipMallocAsync(reinterpret_cast(&p), alloc_size, reinterpret_cast(-1)), + hipErrorInvalidHandle); + } + + SECTION("Size is max size_t") { + HIP_CHECK_ERROR(hipMallocAsync(reinterpret_cast(&p), max_size, stream.stream()), + hipErrorOutOfMemory); + } } diff --git a/projects/hip-tests/catch/unit/memory/hipMallocFromPoolAsync.cc b/projects/hip-tests/catch/unit/memory/hipMallocFromPoolAsync.cc new file mode 100644 index 0000000000..408e43374a --- /dev/null +++ b/projects/hip-tests/catch/unit/memory/hipMallocFromPoolAsync.cc @@ -0,0 +1,149 @@ +/* + 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 + +/** + * @addtogroup hipMallocFromPoolAsync hipMallocFromPoolAsync + * @{ + * @ingroup StreamOTest + * `hipMallocFromPoolAsync(void** dev_ptr, size_t size, hipMemPool_t mem_pool, hipStream_t stream)` + * - Allocates memory from a specified pool with stream ordered semantics + */ + + +/** + * Test Description + * ------------------------ + * - Basic test to verify proper allocation and stream ordering of hipMallocFromPoolAsync when one + * memory allocation is performed. + * Test source + * ------------------------ + * - /unit/memory/hipMallocFromPoolAsync.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipMallocFromPoolAsync_Basic_OneAlloc") { + MallocMemPoolAsync_OneAlloc( + [](void** dev_ptr, size_t size, hipMemPool_t mem_pool, hipStream_t stream) { + return hipMallocFromPoolAsync(dev_ptr, size, mem_pool, stream); + }, + MemPools::created); +} + + +/** + * Test Description + * ------------------------ + * - Basic test to verify proper allocation and stream ordering of hipMallocFromPoolAsync when two + * memory allocations are performed. + * Test source + * ------------------------ + * - /unit/memory/hipMallocFromPoolAsync.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipMallocFromPoolAsync_Basic_TwoAllocs") { + MallocMemPoolAsync_TwoAllocs( + [](void** dev_ptr, size_t size, hipMemPool_t mem_pool, hipStream_t stream) { + return hipMallocFromPoolAsync(dev_ptr, size, mem_pool, stream); + }, + MemPools::created); +} + +/** + * Test Description + * ------------------------ + * - Basic test to verify that memory allocated with hipMallocFromPoolAsync can be properly reused. + * Test source + * ------------------------ + * - /unit/memory/hipMallocFromPoolAsync.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipMallocFromPoolAsync_Basic_Reuse") { + MallocMemPoolAsync_Reuse( + [](void** dev_ptr, size_t size, hipMemPool_t mem_pool, hipStream_t stream) { + return hipMallocFromPoolAsync(dev_ptr, size, mem_pool, stream); + }, + MemPools::created); +} + +/** + * Test Description + * ------------------------ + * - Test to verify hipMallocFromPoolAsync behavior with invalid arguments: + * -# Nullptr dev_ptr + * -# Nullptr mem_pool + * -# Invalid stream handle + * -# Size is max size_t + * + * Test source + * ------------------------ + * - /unit/memory/hipMallocFromPoolAsync.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ +TEST_CASE("Unit_hipMallocFromPoolAsync_Negative_Parameters") { + 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; + } + + void* p = nullptr; + size_t max_size = std::numeric_limits::max(); + size_t alloc_size = 1024; + MemPoolGuard mempool(MemPools::created, device_id); + StreamGuard stream(Streams::created); + + SECTION("dev_ptr is nullptr") { + HIP_CHECK_ERROR(hipMallocFromPoolAsync(nullptr, alloc_size, mempool.mempool(), stream.stream()), + hipErrorInvalidValue); + } + + SECTION("Mempool not created") { + hipMemPool_t dummy_mem_pool = nullptr; + HIP_CHECK_ERROR(hipMallocFromPoolAsync(static_cast(&p), alloc_size, dummy_mem_pool, + stream.stream()), + hipErrorInvalidValue); + } + + SECTION("Invalid stream handle") { + HIP_CHECK_ERROR(hipMallocFromPoolAsync(static_cast(&p), alloc_size, mempool.mempool(), + reinterpret_cast(-1)), + hipErrorInvalidHandle); + } + + SECTION("Size is max size_t") { + HIP_CHECK_ERROR(hipMallocFromPoolAsync(static_cast(&p), max_size, mempool.mempool(), + stream.stream()), + hipErrorOutOfMemory); + } +} diff --git a/projects/hip-tests/catch/unit/memory/mempool_common.hh b/projects/hip-tests/catch/unit/memory/mempool_common.hh index 0f8e6499a8..9bd95f0511 100644 --- a/projects/hip-tests/catch/unit/memory/mempool_common.hh +++ b/projects/hip-tests/catch/unit/memory/mempool_common.hh @@ -19,6 +19,8 @@ #pragma once #include +#include +#include namespace { constexpr hipMemPoolProps kPoolProps = { @@ -28,32 +30,263 @@ constexpr auto wait_ms = 500; } // anonymous namespace -template -__global__ void kernel_500ms(T* host_res, int clk_rate) { +template __global__ void kernel_500ms(T* host_res, int clk_rate) { int tid = threadIdx.x + blockIdx.x * blockDim.x; host_res[tid] = tid + 1; __threadfence_system(); // expecting that the data is getting flushed to host here! - uint64_t start = clock64()/clk_rate, cur; + uint64_t start = clock64() / clk_rate, cur; if (clk_rate > 1) { - do { cur = clock64()/clk_rate-start;}while (cur < wait_ms); + do { + cur = clock64() / clk_rate - start; + } while (cur < wait_ms); } else { - do { cur = clock64()/start;}while (cur < wait_ms); + do { + cur = clock64() / start; + } while (cur < wait_ms); } } -template -__global__ void kernel_500ms_gfx11(T* host_res, int clk_rate) { +template __global__ void kernel_500ms_gfx11(T* host_res, int clk_rate) { #if HT_AMD int tid = threadIdx.x + blockIdx.x * blockDim.x; host_res[tid] = tid + 1; __threadfence_system(); // expecting that the data is getting flushed to host here! - uint64_t start = wall_clock64()/clk_rate, cur; + uint64_t start = wall_clock64() / clk_rate, cur; if (clk_rate > 1) { - do { cur = wall_clock64()/clk_rate-start;}while (cur < wait_ms); + do { + cur = wall_clock64() / clk_rate - start; + } while (cur < wait_ms); } else { - do { cur = wall_clock64()/start;}while (cur < wait_ms); + do { + cur = wall_clock64() / start; + } while (cur < wait_ms); } #endif } + +template void MallocMemPoolAsync_OneAlloc(F malloc_func, const MemPools mempool_type) { + 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 auto allocation_size = GENERATE(kPageSize / 2, kPageSize, kPageSize * 2); + LinearAllocGuard host_alloc(LinearAllocs::hipHostMalloc, allocation_size); + MemPoolGuard mempool(mempool_type, device_id); + + int* alloc_mem; + StreamGuard stream(Streams::created); + + HIP_CHECK(malloc_func(reinterpret_cast(&alloc_mem), allocation_size, mempool.mempool(), + stream.stream())); + + int blocks = 1024; + int clk_rate; + hipMemPoolAttr attr; + if (IsGfx11()) { + HIP_CHECK(hipDeviceGetAttribute(&clk_rate, hipDeviceAttributeWallClockRate, 0)); + kernel_500ms_gfx11<<<32, blocks, 0, stream.stream()>>>(alloc_mem, clk_rate); + } else { + HIP_CHECK(hipDeviceGetAttribute(&clk_rate, hipDeviceAttributeClockRate, 0)); + + kernel_500ms<<<32, blocks, 0, stream.stream()>>>(alloc_mem, clk_rate); + } + + const auto element_count = allocation_size / sizeof(int); + constexpr auto thread_count = 1024; + const auto block_count = element_count / thread_count + 1; + constexpr int expected_value = 17; + VectorSet<<>>(alloc_mem, expected_value, + element_count); + + HIP_CHECK(hipMemcpyAsync(host_alloc.host_ptr(), alloc_mem, allocation_size, hipMemcpyDeviceToHost, + stream.stream())); + + HIP_CHECK(hipFreeAsync(reinterpret_cast(alloc_mem), stream.stream())); + + attr = hipMemPoolAttrReservedMemCurrent; + std::uint64_t res_before_sync = 0; + HIP_CHECK(hipMemPoolGetAttribute(mempool.mempool(), attr, &res_before_sync)); + HIP_CHECK(hipStreamSynchronize(stream.stream())); + + std::uint64_t res_after_sync = 0; + HIP_CHECK(hipMemPoolGetAttribute(mempool.mempool(), attr, &res_after_sync)); + // Sync must release memory to OS + REQUIRE(res_after_sync <= res_before_sync); + + std::uint64_t used_mem = 10; + attr = hipMemPoolAttrUsedMemCurrent; + HIP_CHECK(hipMemPoolGetAttribute(mempool.mempool(), attr, &used_mem)); + REQUIRE(0 == used_mem); + + ArrayFindIfNot(host_alloc.host_ptr(), expected_value, element_count); +} + +template +void MallocMemPoolAsync_TwoAllocs(F malloc_func, const MemPools mempool_type) { + 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 auto allocation_size = GENERATE(kPageSize / 2, kPageSize, kPageSize * 2); + LinearAllocGuard host_alloc(LinearAllocs::hipHostMalloc, allocation_size); + MemPoolGuard mempool(mempool_type, device_id); + + int* alloc_mem1; + int* alloc_mem2; + StreamGuard stream(Streams::created); + + HIP_CHECK(malloc_func(reinterpret_cast(&alloc_mem1), allocation_size, mempool.mempool(), + stream.stream())); + HIP_CHECK(malloc_func(reinterpret_cast(&alloc_mem2), allocation_size, mempool.mempool(), + stream.stream())); + + int blocks = 1024; + int clk_rate; + hipMemPoolAttr attr; + 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); + } + + const auto element_count = allocation_size / sizeof(int); + constexpr auto thread_count = 1024; + const auto block_count = element_count / thread_count + 1; + constexpr int expected_value = 17; + VectorSet<<>>(alloc_mem1, expected_value, + element_count); + HIP_CHECK(hipGetLastError()); + + HIP_CHECK(hipMemcpyAsync(alloc_mem2, alloc_mem1, allocation_size, hipMemcpyDeviceToDevice, + stream.stream())); + + HIP_CHECK(hipMemcpyAsync(host_alloc.host_ptr(), alloc_mem2, allocation_size, + hipMemcpyDeviceToHost, stream.stream())); + + HIP_CHECK(hipFreeAsync(reinterpret_cast(alloc_mem1), stream.stream())); + + attr = hipMemPoolAttrReservedMemCurrent; + std::uint64_t res_before_sync = 0; + HIP_CHECK(hipMemPoolGetAttribute(mempool.mempool(), attr, &res_before_sync)); + HIP_CHECK(hipStreamSynchronize(stream.stream())); + + std::uint64_t res_after_sync = 0; + HIP_CHECK(hipMemPoolGetAttribute(mempool.mempool(), attr, &res_after_sync)); + // Sync must release memory to OS + REQUIRE(res_after_sync <= res_before_sync); + + std::uint64_t used_mem = 0; + attr = hipMemPoolAttrUsedMemCurrent; + HIP_CHECK(hipMemPoolGetAttribute(mempool.mempool(), attr, &used_mem)); + // Make sure the current usage query works - just second buffer is left + REQUIRE(allocation_size == used_mem); + + attr = hipMemPoolAttrUsedMemHigh; + HIP_CHECK(hipMemPoolGetAttribute(mempool.mempool(), attr, &used_mem)); + // Make sure the high watermark usage works - both buffers must be reported + REQUIRE((2 * allocation_size) == used_mem); + + HIP_CHECK(hipFreeAsync(reinterpret_cast(alloc_mem2), stream.stream())); + HIP_CHECK(hipStreamSynchronize(stream.stream())); + + attr = hipMemPoolAttrUsedMemCurrent; + HIP_CHECK(hipMemPoolGetAttribute(mempool.mempool(), attr, &used_mem)); + // Make sure the current usage query works - none of the buffers are used + REQUIRE(0 == used_mem); + + ArrayFindIfNot(host_alloc.host_ptr(), expected_value, element_count); +} + +template void MallocMemPoolAsync_Reuse(F malloc_func, const MemPools mempool_type) { + 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; + } + + MemPoolGuard mempool(mempool_type, device_id); + + int *alloc_mem1, *alloc_mem2, *alloc_mem3; + StreamGuard stream(Streams::created); + + size_t allocation_size1 = kPageSize * kPageSize * 2; + HIP_CHECK(malloc_func(reinterpret_cast(&alloc_mem1), allocation_size1, mempool.mempool(), + stream.stream())); + + size_t allocation_size2 = kPageSize; + HIP_CHECK(malloc_func(reinterpret_cast(&alloc_mem3), 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; + // Not a real free, since kernel isn't done + HIP_CHECK(hipFreeAsync(reinterpret_cast(alloc_mem1), stream.stream())); + + HIP_CHECK(malloc_func(reinterpret_cast(&alloc_mem2), allocation_size1, mempool.mempool(), + stream.stream())); + // Runtime must reuse the pointer + REQUIRE(alloc_mem1 == alloc_mem2); + + // Make a sync before the second kernel launch to make sure memory B isn't gone + HIP_CHECK(hipStreamSynchronize(stream.stream())); + + // Second kernel launch with new memory + if (IsGfx11()) { + kernel_500ms_gfx11<<<32, blocks, 0, stream.stream()>>>(alloc_mem2, clk_rate); + } else { + kernel_500ms<<<32, blocks, 0, stream.stream()>>>(alloc_mem2, clk_rate); + } + + HIP_CHECK(hipStreamSynchronize(stream.stream())); + + attr = hipMemPoolAttrUsedMemCurrent; + std::uint64_t value64 = 0; + HIP_CHECK(hipMemPoolGetAttribute(mempool.mempool(), attr, &value64)); + // Make sure the current usage reports the both buffers + REQUIRE((allocation_size1 + 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())); + attr = hipMemPoolAttrUsedMemCurrent; + HIP_CHECK(hipMemPoolGetAttribute(mempool.mempool(), attr, &value64)); + // Make sure the current usage reports just one buffer, because the above free doesn't hold memory + REQUIRE(allocation_size2 == value64); + + HIP_CHECK(hipFreeAsync(reinterpret_cast(alloc_mem3), stream.stream())); +}