AI/rocm-systems
2
0
Garpu 0
Kode Masalah Tarik Permintaan Actions Packages Projects Rilis Wiki Kegiatan

SWDEV-1 - Merge github PRs to amd-staging

Telusuri Sumber 
- https://github.com/ROCm/hip-tests/pull/154
- https://github.com/ROCm/hip-tests/pull/438
- https://github.com/ROCm/hip-tests/pull/425
- https://github.com/ROCm/hip-tests/pull/424
- https://github.com/ROCm/hip-tests/pull/423
- https://github.com/ROCm/hip-tests/pull/365
- https://github.com/ROCm/hip-tests/pull/356
- https://github.com/ROCm/hip-tests/pull/279
- https://github.com/ROCm/hip-tests/pull/274
- https://github.com/ROCm/hip-tests/pull/190
- https://github.com/ROCm/hip-tests/pull/189
- https://github.com/ROCm/hip-tests/pull/188
- https://github.com/ROCm/hip-tests/pull/156
- https://github.com/ROCm/hip-tests/pull/49
- https://github.com/ROCm/hip-tests/pull/439
- https://github.com/ROCm/hip-tests/pull/437
- https://github.com/ROCm/hip-tests/pull/436
- https://github.com/ROCm/hip-tests/pull/435
- https://github.com/ROCm/hip-tests/pull/193

Change-Id: I2529d0baf0f8d47d6215863321720cde2b1a846c


[ROCm/hip-tests commit: 978e165d33]
This commit is contained in:
Rakesh Roy
2024-02-19 15:14:45 +05:30
orang tua 925af16b13 fcffafb921
melakukan c165105c24
101 mengubah file dengan 11970 tambahan dan 257 penghapusan
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip-tests/catch/unit/memory/CMakeLists.txt
+9 -1
Melihat File
@@ -107,6 +107,12 @@ set(TEST_SRC
hipMemcpyFromSymbol.cc
hipPtrGetAttribute.cc
hipMemPoolApi.cc
hipMemPoolSetGetAccess.cc
hipMemPoolSetGetAttribute.cc
hipMemPoolCreate.cc
hipMemPoolDestroy.cc
hipMemPoolTrimTo.cc
hipMallocFromPoolAsync.cc
hipMemcpyPeer.cc
hipMemcpyPeer_old.cc
hipMemcpyPeerAsync.cc
@@ -151,7 +157,9 @@ set(TEST_SRC
hipStreamAttachMemAsync.cc
hipMemRangeGetAttributes_old.cc
hipMemGetAddressRange.cc
hipArrayGetDescriptor.cc)
hipArrayGetDescriptor.cc
hipMallocMipmappedArray.cc
hipFreeMipmappedArray.cc)
set(NOT_FOR_MI200_AND_ABOVE_TEST hipMallocArray.cc hipArrayCreate.cc) # tests not for MI200+
set(MI200_AND_ABOVE_TARGETS gfx90a gfx940 gfx941 gfx942)
projects/hip-tests/catch/unit/memory/hipFreeAsync.cc
+50 -15
Melihat 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/hipFreeMipmappedArray.cc
+126
Melihat File
@@ -0,0 +1,126 @@
/*
Copyright (c) 2022 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 "hipArrayCommon.hh"
#include "utils.hh"
/*
* hipFreeMipmappedArray API test scenarios
* 1. Check that hipFreeMipmappedArray implicitly synchronises the device.
* 2. Perform multiple allocations and then call hipFreeMipmappedArray on each pointer concurrently (from unique
* threads) for different memory types and different allocation sizes.
* 3. Pass nullptr as argument and check that correct error code is returned.
* 4. Call hipFreeMipmappedArray twice on the same pointer and check that the implementation handles the second
* call correctly.
*/
TEMPLATE_TEST_CASE("Unit_hipFreeMipmappedArrayImplicitSyncArray", "", char, float) {
hipMipmappedArray_t arrayPtr{};
hipExtent extent{};
hipChannelFormatDesc desc = hipCreateChannelDesc<TestType>();
#if HT_AMD
const unsigned int flags = hipArrayDefault;
#else
const unsigned int flags = GENERATE(hipArrayDefault, hipArraySurfaceLoadStore);
#endif
extent.width = GENERATE(64, 256, 1024);
extent.height = GENERATE(64, 256, 1024);
extent.depth = GENERATE(0, 64, 256, 1024);
const unsigned int numLevels = GENERATE(1, 5, 7);
HIP_CHECK(hipMallocMipmappedArray(&arrayPtr, &desc, extent, numLevels, flags));
LaunchDelayKernel(std::chrono::milliseconds{50}, nullptr);
// make sure device is busy
HIP_CHECK_ERROR(hipStreamQuery(nullptr), hipErrorNotReady);
HIP_CHECK(hipFreeMipmappedArray(arrayPtr));
HIP_CHECK(hipStreamQuery(nullptr));
}
TEST_CASE("Unit_hipFreeMipmappedArray_Negative_Nullptr") {
HIP_CHECK_ERROR(hipFreeMipmappedArray(nullptr), hipErrorInvalidValue);
}
TEST_CASE("Unit_hipFreeMipmappedArray_Negative_DoubleFree") {
hipMipmappedArray_t arrayPtr{};
hipExtent extent{};
hipChannelFormatDesc desc = hipCreateChannelDesc<char>();
#if HT_AMD
const unsigned int flags = hipArrayDefault;
#else
const unsigned int flags = GENERATE(hipArrayDefault, hipArraySurfaceLoadStore);
#endif
extent.width = GENERATE(64, 512, 1024);
extent.height = GENERATE(64, 512, 1024);
extent.depth = GENERATE(0, 64, 512, 1024);
const unsigned int numLevels = GENERATE(1, 5, 7);
HIP_CHECK(hipMallocMipmappedArray(&arrayPtr, &desc, extent, numLevels, flags));
HIP_CHECK(hipFreeMipmappedArray(arrayPtr));
HIP_CHECK_ERROR(hipFreeMipmappedArray(arrayPtr), hipErrorContextIsDestroyed);
}
TEMPLATE_TEST_CASE("Unit_hipFreeMipmappedArrayMultiTArray", "", char, int) {
constexpr size_t numAllocs = 10;
std::vector<std::thread> threads;
std::vector<hipMipmappedArray_t> ptrs(numAllocs);
hipExtent extent{};
hipChannelFormatDesc desc = hipCreateChannelDesc<TestType>();
const unsigned int numLevels = GENERATE(1, 5, 7);
#if HT_AMD
const unsigned int flags = hipArrayDefault;
#else
const unsigned int flags = GENERATE(hipArrayDefault, hipArraySurfaceLoadStore);
#endif
extent.width = GENERATE(64, 256, 1024);
extent.height = GENERATE(64, 256, 1024);
extent.depth = GENERATE(0, 64, 256, 1024);
for (auto& ptr : ptrs) {
HIP_CHECK(hipMallocMipmappedArray(&ptr, &desc, extent, numLevels, flags));
}
for (auto ptr : ptrs) {
threads.emplace_back(([ptr] {
HIP_CHECK_THREAD(hipFreeMipmappedArray(ptr));
HIP_CHECK_THREAD(hipStreamQuery(nullptr));
}));
}
for (auto& t : threads) {
t.join();
}
HIP_CHECK_THREAD_FINALIZE();
}
projects/hip-tests/catch/unit/memory/hipMalloc3D.cc
+11
Melihat File
@@ -69,7 +69,18 @@ TEST_CASE("Unit_hipMalloc3D_Basic") {
size_t height{SMALL_SIZE}, depth{SMALL_SIZE};
hipPitchedPtr devPitchedPtr;
hipExtent extent = make_hipExtent(width, height, depth);
size_t tot, avail, ptot, pavail;
HIP_CHECK(hipMemGetInfo(&pavail, &ptot));
REQUIRE(hipMalloc3D(&devPitchedPtr, extent) == hipSuccess);
HIPCHECK(hipFree(devPitchedPtr.ptr));
HIP_CHECK(hipMemGetInfo(&avail, &tot));
if (pavail != avail) {
WARN("Memory leak of hipMalloc3D API in multithreaded scenario");
REQUIRE(false);
}
}
/*
projects/hip-tests/catch/unit/memory/hipMallocAsync.cc
+113 -22
Melihat 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
Melihat 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/hipMallocMipmappedArray.cc
+417
Melihat File
@@ -0,0 +1,417 @@
/*
Copyright (c) 2022 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.
*/
/*
hipMallocMipmappedArray API test scenarios
1. Basic Functionality
2. Negative Scenarios
3. Allocating Small and big chunk data
4. Multithreaded scenario
*/
#include <array>
#include <hip_test_common.hh>
#include "hipArrayCommon.hh"
static constexpr auto ARRAY_SIZE{4};
static constexpr auto BIG_ARRAY_SIZE{100};
static constexpr auto ARRAY_LOOP{100};
/*
* This API verifies memory allocations for small and
* bigger chunks of data.
* Two scenarios are verified in this API
* 1. SmallArray: Allocates ARRAY_SIZE in a loop and
* releases the memory and verifies the meminfo.
* 2. BigArray: Allocates BIG_ARRAY_SIZE in a loop and
* releases the memory and verifies the meminfo
*
* In both cases, the memory info before allocation and
* after releasing the memory should be the same
*
*/
static void MallocMipmappedArray_DiffSizes(int gpu) {
HIP_CHECK_THREAD(hipSetDevice(gpu));
// Use of GENERATE in thead function causes random failures with multithread condition.
std::vector<size_t> runs{ARRAY_SIZE, BIG_ARRAY_SIZE};
for (const auto& size : runs) {
auto numLevelsLimit = floor(log2(size));
for (unsigned int numLevels = 0; numLevels < numLevelsLimit; numLevels++) {
size_t width{size}, height{size}, depth{size};
hipChannelFormatDesc channelDesc = hipCreateChannelDesc<float>();
std::array<hipMipmappedArray_t, ARRAY_LOOP> arr;
size_t pavail, avail, total;
HIP_CHECK_THREAD(hipMemGetInfo(&pavail, &total));
for (int i = 0; i < ARRAY_LOOP; i++) {
HIP_CHECK_THREAD(hipMallocMipmappedArray(&arr[i], &channelDesc,
make_hipExtent(width, height, depth),
(1 + numLevels), hipArrayDefault));
}
for (int i = 0; i < ARRAY_LOOP; i++) {
HIP_CHECK_THREAD(hipFreeMipmappedArray(arr[i]));
}
HIP_CHECK_THREAD(hipMemGetInfo(&avail, &total));
REQUIRE_THREAD(pavail == avail);
}
}
}
TEST_CASE("Unit_hipMallocMipmappedArray_DiffSizes") {
MallocMipmappedArray_DiffSizes(0);
HIP_CHECK_THREAD_FINALIZE();
}
/*
This testcase verifies the hipMallocMipmappedArray API in multithreaded
scenario by launching threads in parallel on multiple GPUs
and verifies the hipMallocMipmappedArray API with small and big chunks data
*/
TEST_CASE("Unit_hipMallocMipmappedArray_MultiThread") {
std::vector<std::thread> threadlist;
int devCnt = 0;
devCnt = HipTest::getDeviceCount();
for (int i = 0; i < devCnt; i++) {
threadlist.push_back(std::thread(MallocMipmappedArray_DiffSizes, i));
}
for (auto& t : threadlist) {
t.join();
}
HIP_CHECK_THREAD_FINALIZE();
}
namespace {
void checkMipmappedArrayIsExpected(hipArray_t level_array,
const hipChannelFormatDesc& expected_desc,
const hipExtent& expected_extent,
const unsigned int expected_flags) {
// hipArrayGetInfo doesn't currently exist (EXSWCPHIPT-87)
#if HT_AMD
std::ignore = level_array;
std::ignore = expected_desc;
std::ignore = expected_extent;
std::ignore = expected_flags;
#else
cudaChannelFormatDesc queried_desc;
cudaExtent queried_extent;
unsigned int queried_flags;
cudaArrayGetInfo(&queried_desc, &queried_extent, &queried_flags, level_array);
REQUIRE(expected_desc.x == queried_desc.x);
REQUIRE(expected_desc.y == queried_desc.y);
REQUIRE(expected_desc.z == queried_desc.z);
REQUIRE(expected_desc.f == queried_desc.f);
REQUIRE(expected_extent.width == queried_extent.width);
REQUIRE(expected_extent.height == queried_extent.height);
REQUIRE(expected_extent.depth == queried_extent.depth);
REQUIRE(expected_flags == queried_flags);
#endif
}
} // namespace
TEMPLATE_TEST_CASE("Unit_hipMallocMipmappedArray_happy", "", char, uint2, int4, short4, float) {
hipMipmappedArray_t array;
const auto desc = hipCreateChannelDesc<TestType>();
#if HT_AMD
const unsigned int flags = hipArrayDefault;
#else
const unsigned int flags =
GENERATE(hipArrayDefault, hipArraySurfaceLoadStore, hipArrayTextureGather);
#endif
constexpr size_t size = 64;
const unsigned int numLevels = GENERATE(1, 3, 5, 7);
std::vector<hipExtent> extents;
extents.reserve(3);
extents.push_back({size, size, 0}); // 2D array
if (flags != hipArrayTextureGather) {
extents.push_back({size, 0, 0}); // 1D array
extents.push_back({size, size, size}); // 3D array
};
for (const auto extent : extents) {
CAPTURE(flags, extent.width, extent.height, extent.depth);
HIP_CHECK(hipMallocMipmappedArray(&array, &desc, extent, numLevels, flags));
hipArray_t hipArray = nullptr;
HIP_CHECK(hipGetMipmappedArrayLevel(&hipArray, array, 0));
checkMipmappedArrayIsExpected(hipArray, desc, extent, flags);
HIP_CHECK(hipFreeMipmappedArray(array));
}
}
#if HT_AMD
constexpr std::array<unsigned int, 1> validFlags{hipArrayDefault};
#else
constexpr std::array<unsigned int, 9> validFlags{
hipArrayDefault,
hipArrayDefault | hipArraySurfaceLoadStore,
hipArrayLayered,
hipArrayLayered | hipArraySurfaceLoadStore,
hipArrayCubemap,
hipArrayCubemap | hipArrayLayered,
hipArrayCubemap | hipArraySurfaceLoadStore,
hipArrayCubemap | hipArrayLayered | hipArraySurfaceLoadStore,
hipArrayTextureGather};
#endif
hipExtent makeMipmappedExtent(unsigned int flag, size_t s) {
if (flag == hipArrayTextureGather) {
return make_hipExtent(s, s, 0);
}
return make_hipExtent(s, s, s);
}
// Providing the array pointer as nullptr should return an error
TEST_CASE("Unit_hipMallocMipmappedArray_Negative_NullArrayPtr") {
hipChannelFormatDesc desc = hipCreateChannelDesc<float4>();
unsigned int numLevels = 1;
constexpr size_t s = 6;
const auto flag = GENERATE(from_range(std::begin(validFlags), std::end(validFlags)));
HIP_CHECK_ERROR(
hipMallocMipmappedArray(nullptr, &desc, makeMipmappedExtent(flag, s), numLevels, flag),
hipErrorInvalidValue);
}
// Providing the description pointer as nullptr should return an error
TEST_CASE("Unit_hipMallocMipmappedArray_Negative_NullDescPtr") {
constexpr size_t s = 6; // 6 to keep cubemap happy
unsigned int numLevels = 1;
hipMipmappedArray_t array;
const auto flag = GENERATE(from_range(std::begin(validFlags), std::end(validFlags)));
HIP_CHECK_ERROR(
hipMallocMipmappedArray(&array, nullptr, makeMipmappedExtent(flag, s), numLevels, flag),
hipErrorInvalidValue);
}
// Zero width arrays are not allowed
TEST_CASE("Unit_hipMallocMipmappedArray_Negative_ZeroWidth") {
constexpr size_t s = 6; // 6 to keep cubemap happy
unsigned int numLevels = 1;
hipMipmappedArray_t array;
hipChannelFormatDesc desc = hipCreateChannelDesc<float4>();
const auto flag = GENERATE(from_range(std::begin(validFlags), std::end(validFlags)));
HIP_CHECK_ERROR(hipMallocMipmappedArray(&array, &desc, make_hipExtent(0, s, s), numLevels, flag),
hipErrorInvalidValue);
}
// Zero height arrays are only allowed for 1D arrays and layered arrays
TEST_CASE("Unit_hipMallocMipmappedArray_Negative_ZeroHeight") {
constexpr size_t s = 6; // 6 to keep cubemap happy
unsigned int numLevels = 1;
hipMipmappedArray_t array;
hipChannelFormatDesc desc = hipCreateChannelDesc<float4>();
std::array<unsigned int, 2> exceptions{hipArrayLayered,
hipArrayLayered | hipArraySurfaceLoadStore};
const auto flag = GENERATE(from_range(std::begin(validFlags), std::end(validFlags)));
if (std::find(std::begin(exceptions), std::end(exceptions), flag) == std::end(exceptions)) {
// flag is not in list of exceptions
HIP_CHECK_ERROR(
hipMallocMipmappedArray(&array, &desc, make_hipExtent(s, 0, s), numLevels, flag),
hipErrorInvalidValue);
}
}
TEST_CASE("Unit_hipMallocMipmappedArray_Negative_InvalidFlags") {
constexpr size_t s = 6; // 6 to keep cubemap happy
unsigned int numLevels = 1;
hipMipmappedArray_t array;
hipChannelFormatDesc desc = hipCreateChannelDesc<float4>();
#if HT_AMD
const unsigned int flag = 0xDEADBEEF;
#else
const unsigned int flag =
GENERATE(0xDEADBEEF, hipArrayTextureGather | hipArraySurfaceLoadStore,
hipArrayTextureGather | hipArrayCubemap,
hipArrayTextureGather | hipArraySurfaceLoadStore | hipArrayCubemap);
#endif
CAPTURE(flag);
REQUIRE(std::find(std::begin(validFlags), std::end(validFlags), flag) == std::end(validFlags));
HIP_CHECK_ERROR(
hipMallocMipmappedArray(&array, &desc, makeMipmappedExtent(flag, s), numLevels, flag),
hipErrorInvalidValue);
}
void testInvalidDescriptionMipmapped(hipChannelFormatDesc desc) {
constexpr size_t s = 6; // 6 to keep cubemap happy
unsigned int numLevels = 1;
hipMipmappedArray_t array;
#if HT_NVIDIA
hipError_t expectedError = hipErrorUnknown;
#else
hipError_t expectedError = hipErrorInvalidValue;
#endif
const auto flag = GENERATE(from_range(std::begin(validFlags), std::end(validFlags)));
HIP_CHECK_ERROR(
hipMallocMipmappedArray(&array, &desc, makeMipmappedExtent(flag, s), numLevels, flag),
expectedError);
}
TEST_CASE("Unit_hipMallocMipmappedArray_Negative_InvalidFormat") {
hipChannelFormatDesc desc = hipCreateChannelDesc<float4>();
desc.f = GENERATE(hipChannelFormatKindNone, 0xBEEF);
testInvalidDescriptionMipmapped(desc);
}
TEST_CASE("Unit_hipMallocMipmappedArray_Negative_BadChannelLayout") {
const int bits = GENERATE(8, 16, 32);
const hipChannelFormatKind formatKind =
GENERATE(hipChannelFormatKindSigned, hipChannelFormatKindUnsigned, hipChannelFormatKindFloat);
if (bits == 8 && formatKind == hipChannelFormatKindFloat) return;
hipChannelFormatDesc desc = GENERATE_COPY(hipCreateChannelDesc(bits, bits, bits, 0, formatKind),
hipCreateChannelDesc(0, bits, bits, 0, formatKind),
hipCreateChannelDesc(0, bits, bits, bits, formatKind),
hipCreateChannelDesc(bits, 0, bits, 0, formatKind),
hipCreateChannelDesc(bits, bits, 0, bits, formatKind),
hipCreateChannelDesc(0, 0, bits, 0, formatKind),
hipCreateChannelDesc(0, 0, bits, bits, formatKind));
INFO("kind: " << channelFormatString(formatKind));
INFO("x: " << desc.x << ", y: " << desc.y << ", z: " << desc.z << ", w: " << desc.w);
testInvalidDescriptionMipmapped(desc);
}
TEST_CASE("Unit_hipMallocMipmappedArray_Negative_8BitFloat") {
hipChannelFormatDesc desc = GENERATE(hipCreateChannelDesc(8, 0, 0, 0, hipChannelFormatKindFloat),
hipCreateChannelDesc(8, 8, 0, 0, hipChannelFormatKindFloat),
hipCreateChannelDesc(8, 8, 8, 8, hipChannelFormatKindFloat));
testInvalidDescriptionMipmapped(desc);
}
TEST_CASE("Unit_hipMallocMipmappedArray_Negative_DifferentChannelSizes") {
const int bitsX = GENERATE(8, 16, 32);
const int bitsY = GENERATE(8, 16, 32);
const int bitsZ = GENERATE(8, 16, 32);
const int bitsW = GENERATE(8, 16, 32);
if (bitsX == bitsY && bitsY == bitsZ && bitsZ == bitsW) return; // skip when they are equal
const hipChannelFormatKind channelFormat =
GENERATE(hipChannelFormatKindSigned, hipChannelFormatKindUnsigned, hipChannelFormatKindFloat);
if (channelFormat == hipChannelFormatKindFloat &&
(bitsX == 8 || bitsY == 8 || bitsZ == 8 || bitsW == 8))
return; // 8 bit floats aren't allowed
hipChannelFormatDesc desc = hipCreateChannelDesc(bitsX, bitsY, bitsZ, bitsW, channelFormat);
INFO("format: " << channelFormatString(channelFormat) << ", x bits: " << bitsX
<< ", y bits: " << bitsY << ", z bits: " << bitsZ << ", w bits: " << bitsW);
testInvalidDescriptionMipmapped(desc);
}
TEST_CASE("Unit_hipMallocMipmappedArray_Negative_BadChannelSize") {
const int badBits = GENERATE(-1, 0, 10, 100);
const hipChannelFormatKind formatKind =
GENERATE(hipChannelFormatKindSigned, hipChannelFormatKindUnsigned, hipChannelFormatKindFloat);
hipChannelFormatDesc desc = hipCreateChannelDesc(badBits, badBits, badBits, badBits, formatKind);
INFO("Number of bits: " << badBits);
testInvalidDescriptionMipmapped(desc);
}
// hipMallocMipmappedArray should handle the max numeric value gracefully.
TEST_CASE("Unit_hipMallocMipmappedArray_Negative_NumericLimit") {
hipMipmappedArray_t arrayPtr;
unsigned int numLevels = 1;
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
size_t size = std::numeric_limits<size_t>::max();
const auto flag = GENERATE(from_range(std::begin(validFlags), std::end(validFlags)));
HIP_CHECK_ERROR(
hipMallocMipmappedArray(&arrayPtr, &desc, makeMipmappedExtent(flag, size), numLevels, flag),
hipErrorInvalidValue);
}
// texture gather arrays are only allowed to be 2D
TEMPLATE_TEST_CASE("Unit_hipMallocMipmappedArray_Negative_Non2DTextureGather", "", char, uchar2,
float2) {
#if HT_AMD
HipTest::HIP_SKIP_TEST("Texture Gather arrays not supported using AMD backend");
return;
#endif
hipMipmappedArray_t array;
unsigned int numLevels = 1;
const auto desc = hipCreateChannelDesc<TestType>();
constexpr unsigned int flags = hipArrayTextureGather;
constexpr size_t size = 64;
const hipExtent extent = GENERATE(make_hipExtent(size, 0, 0), make_hipExtent(size, size, size));
HIP_CHECK_ERROR(hipMallocMipmappedArray(&array, &desc, extent, numLevels, flags),
hipErrorInvalidValue);
}
TEST_CASE("Unit_hipMallocMipmappedArray_Negative_NumLevels") {
hipMipmappedArray_t array;
constexpr size_t size = 6;
unsigned int numLevels = floor(log2(size)) + 2;
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
const auto flag = GENERATE(from_range(std::begin(validFlags), std::end(validFlags)));
HIP_CHECK_ERROR(
hipMallocMipmappedArray(&array, &desc, makeMipmappedExtent(flag, size), numLevels, flag),
hipErrorInvalidValue);
}
TEST_CASE("Unit_hipGetMipmappedArrayLevel_Negative") {
constexpr size_t s = 6;
unsigned int numLevels = 1;
hipMipmappedArray_t array;
hipArray_t level_array;
hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
HIP_CHECK(
hipMallocMipmappedArray(&array, &desc, make_hipExtent(s, s, s), numLevels, hipArrayDefault));
SECTION("Level is invalid") {
HIP_CHECK_ERROR(hipGetMipmappedArrayLevel(&level_array, array, 3), hipErrorInvalidValue);
}
SECTION("Mipmapped array is nullptr") {
HIP_CHECK_ERROR(hipGetMipmappedArrayLevel(&level_array, nullptr, 1),
hipErrorInvalidResourceHandle);
}
HIP_CHECK(hipFreeMipmappedArray(array));
}
projects/hip-tests/catch/unit/memory/hipMemPoolCreate.cc
+93
Melihat File
@@ -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 <hip_test_common.hh>
/**
* @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;
}
}
projects/hip-tests/catch/unit/memory/hipMemPoolDestroy.cc
+71
Melihat File
@@ -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);
}
}
projects/hip-tests/catch/unit/memory/hipMemPoolSetGetAccess.cc
+361
Melihat File
@@ -0,0 +1,361 @@
/*
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 <hip_test_common.hh>
#include <resource_guards.hh>
#include <utils.hh>
/**
* @addtogroup hipMemPoolSetAccess hipMemPoolSetAccess
* @{
* @ingroup StreamOTest
* `hipMemPoolSetAccess(hipMemPool_t mem_pool, const hipMemAccessDesc* desc_list, size_t count)`
* - Controls visibility of the specified pool between devices
*/
__global__ void copyP2PAndScale(int* dst, const int* src, size_t N) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < N) {
// scale & store src vector.
dst[idx] = 2 * src[idx];
}
}
static void MemPoolSetGetAccess(const MemPools mempool_type, int src_device, int dst_device,
hipMemAccessFlags access_flags) {
MemPoolGuard mempool(mempool_type, src_device);
hipMemAccessDesc desc;
memset(&desc, 0, sizeof(hipMemAccessDesc));
desc.location.type = hipMemLocationTypeDevice;
desc.location.id = dst_device;
desc.flags = access_flags;
HIP_CHECK(hipMemPoolSetAccess(mempool.mempool(), &desc, 1));
hipMemAccessFlags flags = hipMemAccessFlagsProtNone;
HIP_CHECK(hipMemPoolGetAccess(&flags, mempool.mempool(), &desc.location));
REQUIRE(flags == access_flags);
}
/**
* Test Description
* ------------------------
* - Basic test to verify hipMemPoolSetAccess/hipMemPoolGetAccess on a single device.
* Test source
* ------------------------
* - /unit/memory/hipMemPoolSetGetAccess.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.0
*/
TEST_CASE("Unit_hipMemPoolSetGetAccess_Positive_Basic") {
const auto device = GENERATE(range(0, HipTest::getDeviceCount()));
int mem_pool_support = 0;
HIP_CHECK(
hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, device));
if (!mem_pool_support) {
SUCCEED("Runtime doesn't support Memory Pool. Skip the test case.");
return;
}
const auto mempool_type = GENERATE(MemPools::dev_default, MemPools::created);
MemPoolSetGetAccess(mempool_type, device, device, hipMemAccessFlagsProtReadWrite);
}
int CheckP2PMemPoolSupport(int src_device, int dst_device) {
int mem_pool_support = 0;
HIP_CHECK(
hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, src_device));
if (mem_pool_support) {
HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported,
dst_device));
}
return mem_pool_support;
}
/**
* Test Description
* ------------------------
* - Basic test to verify hipMemPoolSetAccess/hipMemPoolGetAccess on multiple devices.
* Test source
* ------------------------
* - /unit/memory/hipMemPoolSetGetAccess.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.0
*/
TEST_CASE("Unit_hipMemPoolSetGetAccess_Positive_MultipleGPU") {
const auto device_count = HipTest::getDeviceCount();
if (device_count < 2) {
HipTest::HIP_SKIP_TEST("Skipping because devices < 2");
return;
}
const auto src_device = GENERATE(range(0, HipTest::getDeviceCount()));
const auto dst_device = GENERATE(range(0, HipTest::getDeviceCount()));
INFO("Src device: " << src_device << ", Dst device: " << dst_device);
int mem_pool_support = CheckP2PMemPoolSupport(src_device, dst_device);
if (!mem_pool_support) {
SUCCEED("Runtime doesn't support Memory Pool. Skip the test case.");
return;
}
const auto mempool_type = GENERATE(MemPools::dev_default, MemPools::created);
const auto access_flag = GENERATE(hipMemAccessFlagsProtNone, hipMemAccessFlagsProtRead,
hipMemAccessFlagsProtReadWrite);
int can_access_peer = 0;
HIP_CHECK(hipSetDevice(src_device));
HIP_CHECK(hipDeviceCanAccessPeer(&can_access_peer, src_device, dst_device));
if (can_access_peer) {
MemPoolSetGetAccess(mempool_type, src_device, dst_device, access_flag);
}
}
void MemPoolSetGetAccess_P2P(const MemPools mempool_type) {
const auto src_device = GENERATE(range(0, HipTest::getDeviceCount()));
const auto dst_device = GENERATE(range(0, HipTest::getDeviceCount()));
INFO("Src device: " << src_device << ", Dst device: " << dst_device);
const auto allocation_size = GENERATE(kPageSize / 2, kPageSize, kPageSize * 2);
int mem_pool_support = CheckP2PMemPoolSupport(src_device, dst_device);
if (!mem_pool_support) {
SUCCEED("Runtime doesn't support Memory Pool. Skip the test case.");
return;
}
int *alloc_mem1, *alloc_mem2;
int can_access_peer = 0;
HIP_CHECK(hipSetDevice(src_device));
HIP_CHECK(hipDeviceCanAccessPeer(&can_access_peer, src_device, dst_device));
if (can_access_peer) {
hipEvent_t waitOnStream1;
LinearAllocGuard<int> host_alloc(LinearAllocs::malloc, allocation_size);
HIP_CHECK(hipEventCreate(&waitOnStream1))
StreamGuard stream1(Streams::withFlags, hipStreamNonBlocking);
// Get/create mempool for src_device
MemPoolGuard mempool(mempool_type, src_device);
// Allocate memory in a stream from the pool set above
if (mempool_type == MemPools::dev_default) {
HIP_CHECK(
hipMallocAsync(reinterpret_cast<void**>(&alloc_mem1), allocation_size, stream1.stream()));
} else {
HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast<void**>(&alloc_mem1), allocation_size,
mempool.mempool(), stream1.stream()));
}
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 = 15;
VectorSet<<<block_count, thread_count, 0, stream1.stream()>>>(alloc_mem1, expected_value,
element_count);
HIP_CHECK(hipEventRecord(waitOnStream1, stream1.stream()));
HIP_CHECK(hipSetDevice(dst_device));
StreamGuard stream2(Streams::withFlags, hipStreamNonBlocking);
// Allocate memory in dst device
HIP_CHECK(
hipMallocAsync(reinterpret_cast<void**>(&alloc_mem2), allocation_size, stream2.stream()));
// Setup peer mappings for dst device
hipMemAccessDesc desc;
memset(&desc, 0, sizeof(hipMemAccessDesc));
desc.location.type = hipMemLocationTypeDevice;
desc.location.id = dst_device;
desc.flags = hipMemAccessFlagsProtReadWrite;
HIP_CHECK(hipMemPoolSetAccess(mempool.mempool(), &desc, 1));
hipMemAccessFlags flags = hipMemAccessFlagsProtNone;
HIP_CHECK(hipMemPoolGetAccess(&flags, mempool.mempool(), &desc.location));
REQUIRE(flags == hipMemAccessFlagsProtReadWrite);
HIP_CHECK(hipStreamWaitEvent(stream2.stream(), waitOnStream1, 0));
copyP2PAndScale<<<block_count, thread_count, 0, stream2.stream()>>>(alloc_mem2, alloc_mem1,
element_count);
HIP_CHECK(hipMemcpyAsync(host_alloc.host_ptr(), alloc_mem2, allocation_size,
hipMemcpyDeviceToHost, stream2.stream()));
HIP_CHECK(hipFreeAsync(alloc_mem1, stream2.stream()));
HIP_CHECK(hipFreeAsync(alloc_mem2, stream2.stream()));
HIP_CHECK(hipStreamSynchronize(stream2.stream()));
ArrayFindIfNot(host_alloc.host_ptr(), 2 * expected_value, element_count);
}
}
/**
* Test Description
* ------------------------
* - Basic test to verify peer-to-peer access of stream ordered memory with hipMemPoolSetAccess.
* Test source
* ------------------------
* - /unit/memory/hipMemPoolSetGetAccess.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.0
*/
TEST_CASE("Unit_hipMemPoolSetGetAccess_Positive_P2P") {
const auto device_count = HipTest::getDeviceCount();
if (device_count < 2) {
HipTest::HIP_SKIP_TEST("Skipping because devices < 2");
return;
}
SECTION("Default MemPool") { MemPoolSetGetAccess_P2P(MemPools::dev_default); }
SECTION("Created MemPool") { MemPoolSetGetAccess_P2P(MemPools::created); }
}
/**
* Test Description
* ------------------------
* - Test to verify hipMemPoolSetAccess behavior with invalid arguments:
* -# Nullptr mem_pool
* -# Desc is nullptr and count is > 0
* -# Count > num_device
* -# Invalid desc location type
* -# Invalid desc location id
* -# Revoking access to own memory pool
*
* Test source
* ------------------------
* - /unit/memory/hipMemPoolSetGetAccess.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.0
*/
TEST_CASE("Unit_hipMemPoolSetAccess_Negative_Parameters") {
int device_id = 0;
HIP_CHECK(hipSetDevice(device_id));
MemPoolGuard mempool(MemPools::dev_default, device_id);
int num_dev = 0;
HIP_CHECK(hipGetDeviceCount(&num_dev));
hipMemAccessDesc desc;
memset(&desc, 0, sizeof(hipMemAccessDesc));
desc.location.type = hipMemLocationTypeDevice;
desc.location.id = device_id;
desc.flags = hipMemAccessFlagsProtReadWrite;
SECTION("Mempool is nullptr") {
HIP_CHECK_ERROR(hipMemPoolSetAccess(nullptr, &desc, 1), hipErrorInvalidValue);
}
#if HT_AMD
SECTION("Desc is nullptr and count is > 0") {
HIP_CHECK_ERROR(hipMemPoolSetAccess(mempool.mempool(), nullptr, 1), hipErrorInvalidValue);
}
#endif
SECTION("Count > num_device") {
HIP_CHECK_ERROR(hipMemPoolSetAccess(mempool.mempool(), &desc, (num_dev + 1)),
hipErrorNotSupported);
}
SECTION("Passing invalid desc location type") {
desc.location.type = hipMemLocationTypeInvalid;
HIP_CHECK_ERROR(hipMemPoolSetAccess(mempool.mempool(), &desc, 1), hipErrorNotSupported);
desc.location.type = hipMemLocationTypeDevice;
}
SECTION("Passing invalid desc location id") {
desc.location.id = num_dev;
HIP_CHECK_ERROR(hipMemPoolSetAccess(mempool.mempool(), &desc, 1), hipErrorInvalidDevice);
desc.location.id = device_id;
}
SECTION("Revoking access to own memory pool") {
desc.flags = hipMemAccessFlagsProtNone;
HIP_CHECK_ERROR(hipMemPoolSetAccess(mempool.mempool(), &desc, 1), hipErrorInvalidDevice);
desc.flags = hipMemAccessFlagsProtReadWrite;
}
}
/**
* End doxygen group hipMemPoolSetAccess.
* @}
*/
/**
* @addtogroup hipMemPoolGetAccess hipMemPoolGetAccess
* @{
* @ingroup StreamOTest
* `hipMemPoolGetAccess(hipMemAccessFlags* flags, hipMemPool_t mem_pool, hipMemLocation* location)`
* - Returns the accessibility of a pool from a device
*/
/**
* Test Description
* ------------------------
* - Test to verify hipMemPoolGetAccess behavior with invalid arguments:
* -# Nullptr mem_pool
* -# Flags is nullptr
* -# Invalid location type
* -# Invalid location id
*
* Test source
* ------------------------
* - /unit/memory/hipMemPoolSetGetAccess.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.0
*/
TEST_CASE("Unit_hipMemPoolGetAccess_Negative_Parameters") {
int device_id = 0;
HIP_CHECK(hipSetDevice(device_id));
MemPoolGuard mempool(MemPools::dev_default, device_id);
int num_dev = 0;
HIP_CHECK(hipGetDeviceCount(&num_dev));
hipMemAccessFlags flags = hipMemAccessFlagsProtNone;
hipMemLocation location = {hipMemLocationTypeDevice, device_id};
SECTION("Mempool is nullptr") {
HIP_CHECK_ERROR(hipMemPoolGetAccess(&flags, nullptr, &location), hipErrorInvalidValue);
}
#if HT_AMD
SECTION("Flags is nullptr") {
HIP_CHECK_ERROR(hipMemPoolGetAccess(nullptr, mempool.mempool(), &location),
hipErrorInvalidValue);
}
#endif
SECTION("Passing invalid location type") {
location.type = hipMemLocationTypeInvalid;
HIP_CHECK_ERROR(hipMemPoolGetAccess(&flags, mempool.mempool(), &location),
hipErrorInvalidValue);
location.type = hipMemLocationTypeDevice;
}
SECTION("Passing invalid location id") {
location.id = num_dev;
HIP_CHECK_ERROR(hipMemPoolGetAccess(&flags, mempool.mempool(), &location),
hipErrorInvalidValue);
location.id = device_id;
}
}
projects/hip-tests/catch/unit/memory/hipMemPoolSetGetAttribute.cc
+590
Melihat File
@@ -0,0 +1,590 @@
/*
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 <resource_guards.hh>
#include <utils.hh>
/**
* @addtogroup hipMemPoolSetAttribute hipMemPoolSetAttribute
* @{
* @ingroup StreamOTest
* `hipMemPoolSetAttribute(hipMemPool_t mem_pool, hipMemPoolAttr attr, void* value)`
* - Sets attributes of a memory pool
*/
template <typename T>
static void MemPoolSetGetAttribute(const hipMemPool_t mempool, const hipMemPoolAttr attr,
T& set_value) {
T get_value = 100;
HIP_CHECK(hipMemPoolSetAttribute(mempool, attr, &set_value));
HIP_CHECK(hipMemPoolGetAttribute(mempool, attr, &get_value));
REQUIRE(get_value == set_value);
}
/**
* Test Description
* ------------------------
* - Basic test to verify that default attribute values are correct.
* Test source
* ------------------------
* - /unit/memory/hipMemPoolSetGetAttribute.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.0
*/
TEST_CASE("Unit_hipMemPoolSetGetAttribute_Positive_Default") {
const auto device = GENERATE(range(0, HipTest::getDeviceCount()));
int mem_pool_support = 0;
HIP_CHECK(
hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, device));
if (!mem_pool_support) {
SUCCEED("Runtime doesn't support Memory Pool. Skip the test case.");
return;
}
const auto mempool_type = GENERATE(MemPools::dev_default, MemPools::created);
MemPoolGuard mempool(mempool_type, device);
const auto attr_type =
GENERATE(hipMemPoolReuseFollowEventDependencies, hipMemPoolReuseAllowOpportunistic,
hipMemPoolReuseAllowInternalDependencies);
// Check default value
int def_value = 0;
HIP_CHECK(hipMemPoolGetAttribute(mempool.mempool(), attr_type, &def_value));
REQUIRE(def_value == 1);
// Check if attribute can be disabled
int set_value = 0;
MemPoolSetGetAttribute(mempool.mempool(), attr_type, set_value);
}
/**
* Test Description
* ------------------------
* - Basic test to verify hipMemPoolSetAttribute/hipMemPoolGetAttribute functionality.
* Test source
* ------------------------
* - /unit/memory/hipMemPoolSetGetAttribute.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.0
*/
TEST_CASE("Unit_hipMemPoolSetGetAttribute_Positive_MemBasic") {
const auto device = GENERATE(range(0, HipTest::getDeviceCount()));
int mem_pool_support = 0;
HIP_CHECK(
hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, device));
if (!mem_pool_support) {
SUCCEED("Runtime doesn't support Memory Pool. Skip the test case.");
return;
}
const auto mempool_type = GENERATE(MemPools::dev_default, MemPools::created);
MemPoolGuard mempool(mempool_type, device);
// Check hipMemPoolAttrReleaseThreshold default value
hipMemPoolAttr attr = hipMemPoolAttrReleaseThreshold;
std::uint64_t value64 = 100;
HIP_CHECK(hipMemPoolGetAttribute(mempool.mempool(), attr, &value64));
REQUIRE(value64 == 0);
// Check setting hipMemPoolAttrReleaseThreshold to a value
std::uint64_t set_value64 = kPageSize;
MemPoolSetGetAttribute(mempool.mempool(), hipMemPoolAttrReleaseThreshold, set_value64);
// Check reset of hipMemPoolAttrReservedMemHigh and hipMemPoolAttrUsedMemHigh
set_value64 = 0;
MemPoolSetGetAttribute(mempool.mempool(), hipMemPoolAttrReservedMemHigh, set_value64);
MemPoolSetGetAttribute(mempool.mempool(), hipMemPoolAttrUsedMemHigh, set_value64);
}
/**
* Test Description
* ------------------------
* - Basic test to verify correct behavior of the Opportunistic attribute.
* Test source
* ------------------------
* - /unit/memory/hipMemPoolSetGetAttribute.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.0
*/
TEST_CASE("Unit_hipMemPoolSetAttribute_Opportunistic") {
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(MemPools::created, device_id);
hipMemPoolAttr attr;
int blocks = 2;
int clk_rate;
if (IsGfx11()) {
HIPCHECK(hipDeviceGetAttribute(&clk_rate, hipDeviceAttributeWallClockRate, 0));
} else {
HIPCHECK(hipDeviceGetAttribute(&clk_rate, hipDeviceAttributeClockRate, 0));
}
int *alloc_mem1, *alloc_mem2, *alloc_mem3;
// Create 2 async non-blocking streams
StreamGuard stream1(Streams::withFlags, hipStreamNonBlocking);
StreamGuard stream2(Streams::withFlags, hipStreamNonBlocking);
size_t allocation_size = kPageSize;
HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast<void**>(&alloc_mem3), allocation_size,
mempool.mempool(), stream1.stream()));
int value = 0;
SECTION("Disallow Opportunistic - No Reuse") {
allocation_size = kPageSize * kPageSize * 2;
HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast<void**>(&alloc_mem1), allocation_size,
mempool.mempool(), stream1.stream()));
// Disable all default pool states
attr = hipMemPoolReuseFollowEventDependencies;
HIP_CHECK(hipMemPoolSetAttribute(mempool.mempool(), attr, &value));
attr = hipMemPoolReuseAllowOpportunistic;
HIP_CHECK(hipMemPoolSetAttribute(mempool.mempool(), attr, &value));
attr = hipMemPoolReuseAllowInternalDependencies;
HIP_CHECK(hipMemPoolSetAttribute(mempool.mempool(), attr, &value));
// Run kernel for 500 ms in the first stream
if (IsGfx11()) {
kernel_500ms_gfx11<<<32, blocks, 0, stream1.stream()>>>(alloc_mem1, clk_rate);
} else {
kernel_500ms<<<32, blocks, 0, stream1.stream()>>>(alloc_mem1, clk_rate);
}
// Not a real free, since kernel isn't done
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(alloc_mem1), stream1.stream()));
// Sleep for 1 second GPU should be idle by now
std::this_thread::sleep_for(std::chrono::milliseconds(1000));
// Allocate memory for the second stream
HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast<void**>(&alloc_mem2), allocation_size,
mempool.mempool(), stream2.stream()));
// Without Opportunistic state runtime must allocate another buffer
REQUIRE(alloc_mem1 != alloc_mem2);
// Run kernel with the new memory in the second stream
if (IsGfx11()) {
kernel_500ms_gfx11<<<32, blocks, 0, stream2.stream()>>>(alloc_mem2, clk_rate);
} else {
kernel_500ms<<<32, blocks, 0, stream2.stream()>>>(alloc_mem2, clk_rate);
}
HIP_CHECK(hipStreamSynchronize(stream1.stream()));
HIP_CHECK(hipStreamSynchronize(stream2.stream()));
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(alloc_mem2), stream2.stream()));
}
SECTION("Disallow Opportunistic - Reuse") {
allocation_size = kPageSize * kPageSize * 2;
HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast<void**>(&alloc_mem1), allocation_size,
mempool.mempool(), stream1.stream()));
// Disable all default pool states
attr = hipMemPoolReuseFollowEventDependencies;
HIP_CHECK(hipMemPoolSetAttribute(mempool.mempool(), attr, &value));
attr = hipMemPoolReuseAllowOpportunistic;
HIP_CHECK(hipMemPoolSetAttribute(mempool.mempool(), attr, &value));
attr = hipMemPoolReuseAllowInternalDependencies;
HIP_CHECK(hipMemPoolSetAttribute(mempool.mempool(), attr, &value));
// Run kernel for 500 ms in the first stream
if (IsGfx11()) {
kernel_500ms_gfx11<<<32, blocks, 0, stream1.stream()>>>(alloc_mem1, clk_rate);
} else {
kernel_500ms<<<32, blocks, 0, stream1.stream()>>>(alloc_mem1, clk_rate);
}
// Not a real free, since kernel isn't done
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(alloc_mem1), stream1.stream()));
// Sleep for 1 second GPU should be idle by now
std::this_thread::sleep_for(std::chrono::milliseconds(1000));
// Allocate memory for the second stream
HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast<void**>(&alloc_mem2), allocation_size,
mempool.mempool(), stream1.stream()));
// Without Opportunistic state runtime must allocate another buffer
REQUIRE(alloc_mem1 == alloc_mem2);
// Run kernel with the new memory in the second stream
if (IsGfx11()) {
kernel_500ms_gfx11<<<32, blocks, 0, stream1.stream()>>>(alloc_mem2, clk_rate);
} else {
kernel_500ms<<<32, blocks, 0, stream1.stream()>>>(alloc_mem2, clk_rate);
}
HIP_CHECK(hipStreamSynchronize(stream1.stream()));
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(alloc_mem2), stream1.stream()));
}
SECTION("Allow Opportunistic - Reuse") {
allocation_size = kPageSize * kPageSize * 2;
HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast<void**>(&alloc_mem1), allocation_size,
mempool.mempool(), stream1.stream()));
value = 1;
attr = hipMemPoolReuseAllowOpportunistic;
// Enable Opportunistic
HIP_CHECK(hipMemPoolSetAttribute(mempool.mempool(), attr, &value));
// Run kernel for 500 ms in the first stream
if (IsGfx11()) {
HIP_CHECK(hipDeviceGetAttribute(&clk_rate, hipDeviceAttributeWallClockRate, 0));
kernel_500ms_gfx11<<<32, blocks, 0, stream1.stream()>>>(alloc_mem1, clk_rate);
} else {
HIP_CHECK(hipDeviceGetAttribute(&clk_rate, hipDeviceAttributeClockRate, 0));
kernel_500ms<<<32, blocks, 0, stream1.stream()>>>(alloc_mem1, clk_rate);
}
// Not a real free, since kernel isn't done
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(alloc_mem1), stream1.stream()));
// Sleep for 1 second GPU should be idle by now
std::this_thread::sleep_for(std::chrono::milliseconds(1000));
// Allocate memory for the second stream
HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast<void**>(&alloc_mem2), allocation_size,
mempool.mempool(), stream2.stream()));
// With Opportunistic state runtime will reuse freed buffer A
REQUIRE(alloc_mem1 == alloc_mem2);
// Run kernel with the new memory in the second stream
if (IsGfx11()) {
kernel_500ms_gfx11<<<32, blocks, 0, stream2.stream()>>>(alloc_mem2, clk_rate);
} else {
kernel_500ms<<<32, blocks, 0, stream2.stream()>>>(alloc_mem2, clk_rate);
}
HIP_CHECK(hipStreamSynchronize(stream1.stream()));
HIP_CHECK(hipStreamSynchronize(stream2.stream()));
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(alloc_mem2), stream2.stream()));
}
SECTION("Allow Opportunistic - No Reuse") {
allocation_size = kPageSize * kPageSize * 2;
HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast<void**>(&alloc_mem1), allocation_size,
mempool.mempool(), stream1.stream()));
value = 1;
attr = hipMemPoolReuseAllowOpportunistic;
// Enable Opportunistic
HIP_CHECK(hipMemPoolSetAttribute(mempool.mempool(), attr, &value));
// Run kernel for 500 ms in the first stream
if (IsGfx11()) {
kernel_500ms_gfx11<<<32, blocks, 0, stream1.stream()>>>(alloc_mem1, clk_rate);
} else {
kernel_500ms<<<32, blocks, 0, stream1.stream()>>>(alloc_mem1, clk_rate);
}
// Not a real free, since kernel isn't done
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(alloc_mem1), stream1.stream()));
// Allocate memory for the second stream
HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast<void**>(&alloc_mem2), allocation_size,
mempool.mempool(), stream2.stream()));
// With Opportunistic state runtime can't reuse freed buffer A, because it's still busy with the
// kernel
REQUIRE(alloc_mem1 != alloc_mem2);
// Run kernel with the new memory in the second stream
if (IsGfx11()) {
kernel_500ms_gfx11<<<32, blocks, 0, stream2.stream()>>>(alloc_mem2, clk_rate);
} else {
kernel_500ms<<<32, blocks, 0, stream2.stream()>>>(alloc_mem2, clk_rate);
}
HIP_CHECK(hipStreamSynchronize(stream1.stream()));
HIP_CHECK(hipStreamSynchronize(stream2.stream()));
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(alloc_mem2), stream2.stream()));
}
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(alloc_mem3), stream1.stream()));
}
/**
* Test Description
* ------------------------
* - Basic test to verify correct behavior of the EventDependencies attribute.
* Test source
* ------------------------
* - /unit/memory/hipMemPoolSetGetAttribute.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.0
*/
TEST_CASE("Unit_hipMemPoolSetAttribute_EventDependencies") {
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(MemPools::created, device_id);
hipMemPoolAttr attr;
int blocks = 2;
int clk_rate;
if (IsGfx11()) {
HIPCHECK(hipDeviceGetAttribute(&clk_rate, hipDeviceAttributeWallClockRate, 0));
} else {
HIPCHECK(hipDeviceGetAttribute(&clk_rate, hipDeviceAttributeClockRate, 0));
}
int *alloc_mem1, *alloc_mem2, *alloc_mem3;
// Create 2 async non-blocking streams
StreamGuard stream1(Streams::withFlags, hipStreamNonBlocking);
StreamGuard stream2(Streams::withFlags, hipStreamNonBlocking);
hipEvent_t event;
HIP_CHECK(hipEventCreate(&event));
size_t allocation_size = kPageSize;
HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast<void**>(&alloc_mem3), allocation_size,
mempool.mempool(), stream1.stream()));
int value = 0;
SECTION("Allow Event Dependencies - Reuse") {
allocation_size = kPageSize * kPageSize * 2;
HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast<void**>(&alloc_mem1), allocation_size,
mempool.mempool(), stream1.stream()));
value = 1;
attr = hipMemPoolReuseFollowEventDependencies;
// Enable Opportunistic
HIP_CHECK(hipMemPoolSetAttribute(mempool.mempool(), attr, &value));
// Run kernel for 500 ms in the first stream
if (IsGfx11()) {
HIP_CHECK(hipDeviceGetAttribute(&clk_rate, hipDeviceAttributeWallClockRate, 0));
kernel_500ms_gfx11<<<32, blocks, 0, stream1.stream()>>>(alloc_mem1, clk_rate);
} else {
HIP_CHECK(hipDeviceGetAttribute(&clk_rate, hipDeviceAttributeClockRate, 0));
kernel_500ms<<<32, blocks, 0, stream1.stream()>>>(alloc_mem1, clk_rate);
}
// Not a real free, since kernel isn't done
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(alloc_mem1), stream1.stream()));
HIP_CHECK(hipEventRecord(event, stream1.stream()));
HIP_CHECK(hipStreamWaitEvent(stream2.stream(), event, 0));
// Allocate memory for the second stream
HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast<void**>(&alloc_mem2), allocation_size,
mempool.mempool(), stream2.stream()));
// With Opportunistic state runtime will reuse freed buffer A
REQUIRE(alloc_mem1 == alloc_mem2);
// Run kernel with the new memory in the second stream
if (IsGfx11()) {
kernel_500ms_gfx11<<<32, blocks, 0, stream2.stream()>>>(alloc_mem2, clk_rate);
} else {
kernel_500ms<<<32, blocks, 0, stream2.stream()>>>(alloc_mem2, clk_rate);
}
HIP_CHECK(hipStreamSynchronize(stream1.stream()));
HIP_CHECK(hipStreamSynchronize(stream2.stream()));
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(alloc_mem2), stream2.stream()));
}
SECTION("Disallow Event Dependencies - No Reuse") {
allocation_size = kPageSize * kPageSize * 2;
HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast<void**>(&alloc_mem1), allocation_size,
mempool.mempool(), stream1.stream()));
value = 0;
attr = hipMemPoolReuseFollowEventDependencies;
// Enable Opportunistic
HIP_CHECK(hipMemPoolSetAttribute(mempool.mempool(), attr, &value));
// Run kernel for 500 ms in the first stream
if (IsGfx11()) {
kernel_500ms_gfx11<<<32, blocks, 0, stream1.stream()>>>(alloc_mem1, clk_rate);
} else {
kernel_500ms<<<32, blocks, 0, stream1.stream()>>>(alloc_mem1, clk_rate);
}
// Not a real free, since kernel isn't done
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(alloc_mem1), stream1.stream()));
HIP_CHECK(hipEventRecord(event, stream1.stream()));
HIP_CHECK(hipStreamWaitEvent(stream2.stream(), event, 0));
// Allocate memory for the second stream
HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast<void**>(&alloc_mem2), allocation_size,
mempool.mempool(), stream2.stream()));
// With Opportunistic state runtime can't reuse freed buffer A, because it's still busy with the
// kernel
REQUIRE(alloc_mem1 != alloc_mem2);
// Run kernel with the new memory in the second stream
if (IsGfx11()) {
kernel_500ms_gfx11<<<32, blocks, 0, stream2.stream()>>>(alloc_mem2, clk_rate);
} else {
kernel_500ms<<<32, blocks, 0, stream2.stream()>>>(alloc_mem2, clk_rate);
}
HIP_CHECK(hipStreamSynchronize(stream1.stream()));
HIP_CHECK(hipStreamSynchronize(stream2.stream()));
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(alloc_mem2), stream2.stream()));
}
HIP_CHECK(hipFreeAsync(reinterpret_cast<void*>(alloc_mem3), stream1.stream()));
HIP_CHECK(hipEventDestroy(event));
}
/**
* Test Description
* ------------------------
* - Test to verify hipMemPoolSetAttribute behavior with invalid arguments:
* -# Nullptr mem_pool
* -# Attribute value is not valid
* -# Nullptr value
* -# hipMemPoolAttrReservedMemHigh set to non-zero
* -# IhipMemPoolAttrUsedMemHigh set to non-zero
*
* Test source
* ------------------------
* - /unit/memory/hipMemPoolSetGetAttribute.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.0
*/
TEST_CASE("Unit_hipMemPoolSetAttribute_Negative_Parameters") {
int device_id = 0;
HIP_CHECK(hipSetDevice(device_id));
MemPoolGuard mempool(MemPools::dev_default, device_id);
hipMemPoolAttr attr = hipMemPoolReuseFollowEventDependencies;
int set_value = 0;
std::uint64_t set_value64 = 0;
SECTION("Mempool is nullptr") {
HIP_CHECK_ERROR(hipMemPoolSetAttribute(nullptr, attr, &set_value), hipErrorInvalidValue);
}
SECTION("Attribute value is not valid") {
HIP_CHECK_ERROR(
hipMemPoolSetAttribute(mempool.mempool(), static_cast<hipMemPoolAttr>(0x9), &set_value),
hipErrorInvalidValue);
}
#if HT_AMD
SECTION("Set values is nullptr") {
HIP_CHECK_ERROR(hipMemPoolSetAttribute(mempool.mempool(), attr, nullptr), hipErrorInvalidValue);
}
#endif
SECTION("Set hipMemPoolAttrReservedMemHigh to non-zero") {
hipMemPoolAttr attr = hipMemPoolAttrReservedMemHigh;
set_value64 = 1;
HIP_CHECK_ERROR((hipMemPoolSetAttribute(mempool.mempool(), attr, &set_value64)),
hipErrorInvalidValue);
}
SECTION("Set hipMemPoolAttrUsedMemHigh to non-zero") {
hipMemPoolAttr attr = hipMemPoolAttrUsedMemHigh;
set_value64 = 1;
HIP_CHECK_ERROR((hipMemPoolSetAttribute(mempool.mempool(), attr, &set_value64)),
hipErrorInvalidValue);
}
}
/**
* End doxygen group hipMemPoolSetAttribute.
* @}
*/
/**
* @addtogroup hipMemPoolGetAttribute hipMemPoolGetAttribute
* @{
* @ingroup StreamOTest
* `hipMemPoolGetAttribute(hipMemPool_t mem_pool, hipMemPoolAttr attr, void* value)`
* - Gets attributes of a memory pool
*/
/**
* Test Description
* ------------------------
* - Test to verify hipMemPoolGetAttribute behavior with invalid arguments:
* -# Nullptr mem_pool
* -# Attribute value is not valid
* -# Nullptr value
*
* Test source
* ------------------------
* - /unit/memory/hipMemPoolSetGetAttribute.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.0
*/
TEST_CASE("Unit_hipMemPoolGetAttribute_Negative_Parameters") {
int device_id = 0;
HIP_CHECK(hipSetDevice(device_id));
MemPoolGuard mempool(MemPools::dev_default, device_id);
hipMemPoolAttr attr = hipMemPoolReuseFollowEventDependencies;
int get_value = 0;
SECTION("Mempool is nullptr") {
HIP_CHECK_ERROR(hipMemPoolGetAttribute(nullptr, attr, &get_value), hipErrorInvalidValue);
}
SECTION("Attribute value is not valid") {
HIP_CHECK_ERROR(
hipMemPoolGetAttribute(mempool.mempool(), static_cast<hipMemPoolAttr>(0x9), &get_value),
hipErrorInvalidValue);
}
SECTION("Get values is nullptr") {
HIP_CHECK_ERROR(hipMemPoolGetAttribute(mempool.mempool(), attr, nullptr), hipErrorInvalidValue);
}
}
projects/hip-tests/catch/unit/memory/hipMemPoolTrimTo.cc
+165
Melihat File
@@ -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 <resource_guards.hh>
#include <utils.hh>
/**
* @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<void**>(&alloc_mem1), allocation_size1,
mempool.mempool(), stream.stream()));
HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast<void**>(&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<void*>(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<void*>(alloc_mem2), stream.stream()));
}
projects/hip-tests/catch/unit/memory/mempool_common.hh
+292
Melihat File
@@ -0,0 +1,292 @@
/*
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.
*/
#pragma once
#include <hip_test_common.hh>
#include <resource_guards.hh>
#include <utils.hh>
namespace {
constexpr hipMemPoolProps kPoolProps = {
hipMemAllocationTypePinned, hipMemHandleTypeNone, {hipMemLocationTypeDevice, 0}, nullptr, {0}};
constexpr auto wait_ms = 500;
} // anonymous namespace
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;
if (clk_rate > 1) {
do {
cur = clock64() / clk_rate - start;
} while (cur < wait_ms);
} else {
do {
cur = clock64() / start;
} while (cur < wait_ms);
}
}
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;
if (clk_rate > 1) {
do {
cur = wall_clock64() / clk_rate - start;
} while (cur < wait_ms);
} else {
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()));
}