AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

EXSWHTEC-35 - Implement tests for hipMallocAsync and hipMallocFromPoolAsync (#438)

Browse Source 
Change-Id: I5cc2ec8078d0de3ddc447ba4bdca74a159f77b12


[ROCm/hip-tests commit: 54f0642d76]
This commit is contained in:
Nives Vukovic
2023-12-15 16:59:04 +05:30
committed by Rakesh Roy
parent f6402bbc10
commit 7c8f7051a3
8 changed files with 564 additions and 51 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip-tests/catch/hipTestMain/config/config_amd_linux
-3
View File
@@ -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 ===",
projects/hip-tests/catch/hipTestMain/config/config_nvidia_linux.json
+1 -1
View File
@@ -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",
projects/hip-tests/catch/include/hip_test_defgroups.hh
+7
View File
@@ -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
* @{
projects/hip-tests/catch/unit/memory/CMakeLists.txt
+1
View File
@@ -112,6 +112,7 @@ set(TEST_SRC
hipMemPoolCreate.cc
hipMemPoolDestroy.cc
hipMemPoolTrimTo.cc
hipMallocFromPoolAsync.cc
hipMemcpyPeer.cc
hipMemcpyPeer_old.cc
hipMemcpyPeerAsync.cc
projects/hip-tests/catch/unit/memory/hipFreeAsync.cc
+50 -15
View File
@@ -21,24 +21,59 @@ THE SOFTWARE.
*/
#include <hip_test_common.hh>
#include <resource_guards.hh>
#include <utils.hh>
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<void**>(&p), 100, stream));
HIP_CHECK(hipStreamSynchronize(stream));
hipError_t error = hipFreeAsync(p, reinterpret_cast<hipStream_t>(-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<void**>(&p), alloc_size, stream.stream()));
HIP_CHECK(hipStreamSynchronize(stream.stream()));
HIP_CHECK_ERROR(hipFreeAsync(p, reinterpret_cast<hipStream_t>(-1)), hipErrorInvalidHandle);
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(p), stream.stream()));
HIP_CHECK(hipStreamSynchronize(stream.stream()));
}
SECTION("Double free") {
HIP_CHECK(hipMallocAsync(reinterpret_cast<void**>(&p), alloc_size, stream.stream()));
HIP_CHECK(hipStreamSynchronize(stream.stream()));
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(p), stream.stream()));
HIP_CHECK(hipStreamSynchronize(stream.stream()));
HIP_CHECK_ERROR(hipFreeAsync(reinterpret_cast<void*>(p), stream.stream()),
hipErrorInvalidValue);
}
}
projects/hip-tests/catch/unit/memory/hipMallocAsync.cc
+113 -22
View File
@@ -17,31 +17,122 @@
THE SOFTWARE.
*/
#include <hip_test_common.hh>
#include <hip_test_checkers.hh>
#include "mempool_common.hh"
#include <limits>
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<size_t>::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<void**>(&p), 100, reinterpret_cast<hipStream_t>(-1)) !=
hipSuccess);
}
SECTION("out of memory") {
REQUIRE(hipMallocAsync(static_cast<void**>(&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<size_t>::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<void**>(&p), alloc_size, reinterpret_cast<hipStream_t>(-1)),
hipErrorInvalidHandle);
}
SECTION("Size is max size_t") {
HIP_CHECK_ERROR(hipMallocAsync(reinterpret_cast<void**>(&p), max_size, stream.stream()),
hipErrorOutOfMemory);
}
}
projects/hip-tests/catch/unit/memory/hipMallocFromPoolAsync.cc
+149
View File
@@ -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 <limits>
/**
* @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<size_t>::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<void**>(&p), alloc_size, dummy_mem_pool,
stream.stream()),
hipErrorInvalidValue);
}
SECTION("Invalid stream handle") {
HIP_CHECK_ERROR(hipMallocFromPoolAsync(static_cast<void**>(&p), alloc_size, mempool.mempool(),
reinterpret_cast<hipStream_t>(-1)),
hipErrorInvalidHandle);
}
SECTION("Size is max size_t") {
HIP_CHECK_ERROR(hipMallocFromPoolAsync(static_cast<void**>(&p), max_size, mempool.mempool(),
stream.stream()),
hipErrorOutOfMemory);
}
}
projects/hip-tests/catch/unit/memory/mempool_common.hh
+243 -10
View File
@@ -19,6 +19,8 @@
#pragma once
#include <hip_test_common.hh>
#include <resource_guards.hh>
#include <utils.hh>
namespace {
constexpr hipMemPoolProps kPoolProps = {
@@ -28,32 +30,263 @@ constexpr auto wait_ms = 500;
} // anonymous namespace
template <typename T>
__global__ void kernel_500ms(T* host_res, int clk_rate) {
template <typename T> __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 <typename T>
__global__ void kernel_500ms_gfx11(T* host_res, int clk_rate) {
template <typename T> __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 <typename F> 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<int> 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<void**>(&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<<<block_count, thread_count, 0, stream.stream()>>>(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<void*>(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 <typename F>
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<int> 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<void**>(&alloc_mem1), allocation_size, mempool.mempool(),
stream.stream()));
HIP_CHECK(malloc_func(reinterpret_cast<void**>(&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<<<block_count, thread_count, 0, stream.stream()>>>(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<void*>(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<void*>(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 <typename F> 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<void**>(&alloc_mem1), allocation_size1, mempool.mempool(),
stream.stream()));
size_t allocation_size2 = kPageSize;
HIP_CHECK(malloc_func(reinterpret_cast<void**>(&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<void*>(alloc_mem1), stream.stream()));
HIP_CHECK(malloc_func(reinterpret_cast<void**>(&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<void*>(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<void*>(alloc_mem3), stream.stream()));
}