EXSWHTEC-69 - Implement tests for hipMemRangeGetAttributes (#51)
- Negative parameter tests for hipMemRangeGetAttribute and hipMemRangeGetAttributes
- Validate the behavior of hipMemRangeGetAttribute for hipMemRangeAttributeReadMostly
- Validate the behavior of hipMemRangeGetAttribute for hipMemRangeAttributePreferredLocation
- Validate the behavior of hipMemRangeGetAttribute for hipMemRangeAttributeLastPrefetchLocation
- Validate the behavior of hipMemRangeGetAttribute for hipMemRangeAttributeAccessedBy
- Validate the behavior of hipMemRangeGetAttributes
[ROCm/hip-tests commit: 203b994230]
Este commit está contenido en:
Mirza Halilčević
GitHub
@@ -196,11 +196,11 @@ template <typename T> class DrvArrayAllocGuard {
|
||||
const hipExtent extent_;
|
||||
};
|
||||
|
||||
enum class Streams { nullstream, perThread, created };
|
||||
enum class Streams { nullstream, perThread, created, withFlags, withPriority };
|
||||
|
||||
class StreamGuard {
|
||||
public:
|
||||
StreamGuard(const Streams stream_type) : stream_type_{stream_type} {
|
||||
StreamGuard(const Streams stream_type, unsigned int flags = hipStreamDefault, int priority = 0) : stream_type_{stream_type}, flags_{flags}, priority_{priority} {
|
||||
switch (stream_type_) {
|
||||
case Streams::nullstream:
|
||||
stream_ = nullptr;
|
||||
@@ -210,6 +210,11 @@ class StreamGuard {
|
||||
break;
|
||||
case Streams::created:
|
||||
HIP_CHECK(hipStreamCreate(&stream_));
|
||||
break;
|
||||
case Streams::withFlags:
|
||||
HIP_CHECK(hipStreamCreateWithFlags(&stream_, flags_));
|
||||
case Streams::withPriority:
|
||||
HIP_CHECK(hipStreamCreateWithPriority(&stream_, flags_, priority_));
|
||||
}
|
||||
}
|
||||
|
||||
@@ -226,5 +231,53 @@ class StreamGuard {
|
||||
|
||||
private:
|
||||
const Streams stream_type_;
|
||||
unsigned int flags_;
|
||||
int priority_;
|
||||
hipStream_t stream_;
|
||||
};
|
||||
|
||||
class EventsGuard {
|
||||
public:
|
||||
EventsGuard(size_t N) : events_(N) {
|
||||
for (auto &e : events_) HIP_CHECK(hipEventCreate(&e));
|
||||
}
|
||||
|
||||
EventsGuard(const EventsGuard&) = delete;
|
||||
EventsGuard(EventsGuard&&) = delete;
|
||||
|
||||
~EventsGuard() {
|
||||
for (auto &e : events_) static_cast<void>(hipEventDestroy(e));
|
||||
}
|
||||
|
||||
hipEvent_t& operator[](int index) { return events_[index]; }
|
||||
|
||||
operator hipEvent_t() const { return events_.at(0); }
|
||||
|
||||
std::vector<hipEvent_t>& event_list() { return events_; }
|
||||
|
||||
private:
|
||||
std::vector<hipEvent_t> events_;
|
||||
};
|
||||
|
||||
class StreamsGuard {
|
||||
public:
|
||||
StreamsGuard(size_t N) : streams_(N) {
|
||||
for (auto &s : streams_) HIP_CHECK(hipStreamCreate(&s));
|
||||
}
|
||||
|
||||
StreamsGuard(const StreamsGuard&) = delete;
|
||||
StreamsGuard(StreamsGuard&&) = delete;
|
||||
|
||||
~StreamsGuard() {
|
||||
for (auto &s : streams_) static_cast<void>(hipStreamDestroy(s));
|
||||
}
|
||||
|
||||
hipStream_t& operator[](int index) { return streams_[index]; }
|
||||
|
||||
operator hipStream_t() const { return streams_.at(0); }
|
||||
|
||||
std::vector<hipStream_t>& stream_list() { return streams_; }
|
||||
|
||||
private:
|
||||
std::vector<hipStream_t> streams_;
|
||||
};
|
||||
|
||||
@@ -66,6 +66,7 @@ set(TEST_SRC
|
||||
hipMemCoherencyTst.cc
|
||||
hipMallocManaged.cc
|
||||
hipMemRangeGetAttribute.cc
|
||||
hipMemRangeGetAttribute_old.cc
|
||||
hipMemcpyFromSymbol.cc
|
||||
hipPtrGetAttribute.cc
|
||||
hipMemPoolApi.cc
|
||||
@@ -111,6 +112,7 @@ set(TEST_SRC
|
||||
hipMemsetAsync.cc
|
||||
hipMemAdvise.cc
|
||||
hipMemRangeGetAttributes.cc
|
||||
hipMemRangeGetAttributes_old.cc
|
||||
hipMemGetAddressRange.cc
|
||||
)
|
||||
else()
|
||||
@@ -157,6 +159,7 @@ set(TEST_SRC
|
||||
hipMemAdviseMmap.cc
|
||||
hipMallocManaged.cc
|
||||
hipMemRangeGetAttribute.cc
|
||||
hipMemRangeGetAttribute_old.cc
|
||||
hipMemcpyFromSymbol.cc
|
||||
hipPtrGetAttribute.cc
|
||||
hipMemPoolApi.cc
|
||||
@@ -198,6 +201,7 @@ set(TEST_SRC
|
||||
hipMemsetAsync.cc
|
||||
hipMemAdvise.cc
|
||||
hipMemRangeGetAttributes.cc
|
||||
hipMemRangeGetAttributes_old.cc
|
||||
hipGetSymbolSizeAddress.cc
|
||||
hipMemGetAddressRange.cc
|
||||
)
|
||||
|
||||
@@ -1,408 +1,359 @@
|
||||
/*
|
||||
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
|
||||
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 WARRANNTY OF ANY KIND, EXPRESS OR
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
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 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
|
||||
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.
|
||||
*/
|
||||
|
||||
/* Test Case Description:
|
||||
Scenario-1: The following function tests the count parameter(last param) to
|
||||
hipMemRangeGetAttribute api by passing possible extreme values.
|
||||
Curently the only way to test if count param working properly is to verify
|
||||
the first parameter of hipMemRangeGetAttribute() api has value 1 stored
|
||||
|
||||
Scenario-2: This test case checks the behavior of hipMemRangeGetAttribute() with
|
||||
AccessedBy flag is consistent with cuda's counter part
|
||||
|
||||
Scenario-3: Allocate 4 * page size of memory with the flag hipMemAttachGloal. Advise
|
||||
AccessedBy, ReadMostly and PreferredLocation to first half(2*pageSz) of the
|
||||
memory and probe the for the flags which are set earlier using
|
||||
hipMemRangeGetAttribute() api for the full size(4*PageSz).
|
||||
|
||||
|
||||
Scenario-4: The following scenarios tests that probing the attributes which are not set
|
||||
by hipMemAdvise() but being probed using hipMemRangeGetAttribute() should
|
||||
not result in a crash
|
||||
|
||||
Scenario-5: The following scenario is a simple test which does the following:
|
||||
Allocate Hmm memory --> hipMemPrefetchAsync() to device 0 and then
|
||||
probe LastPrefetchLocation attribute using hipMemRangeGetAttribute
|
||||
|
||||
Scenario-6: The following Test Case does negative tests on hipMemRangeGetAttribute()*/
|
||||
|
||||
#include <hip/hip_runtime_api.h>
|
||||
#include <hip_test_common.hh>
|
||||
#include <stdlib.h>
|
||||
#ifdef __linux__
|
||||
#include <unistd.h>
|
||||
#include <sys/sysinfo.h>
|
||||
#endif
|
||||
#include <resource_guards.hh>
|
||||
#include <utils.hh>
|
||||
|
||||
static bool CheckError(hipError_t err, int LineNo) {
|
||||
if (err == hipSuccess) {
|
||||
WARN("Error expected but received hipSuccess at line no.:"
|
||||
<< LineNo);
|
||||
return false;
|
||||
} else {
|
||||
return true;
|
||||
TEST_CASE("Unit_hipMemRangeGetAttribute_Positive_ReadMostly_Basic") {
|
||||
if (!DeviceAttributesSupport(0, hipDeviceAttributeManagedMemory)) {
|
||||
HipTest::HIP_SKIP_TEST("Managed memory not supported");
|
||||
return;
|
||||
}
|
||||
|
||||
LinearAllocGuard<void> allocation(LinearAllocs::hipMallocManaged, kPageSize);
|
||||
|
||||
int32_t data;
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(data), hipMemRangeAttributeReadMostly,
|
||||
allocation.ptr(), kPageSize));
|
||||
|
||||
REQUIRE(data == 0);
|
||||
|
||||
HIP_CHECK(hipMemAdvise(allocation.ptr(), kPageSize, hipMemAdviseSetReadMostly, 0));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(data), hipMemRangeAttributeReadMostly,
|
||||
allocation.ptr(), kPageSize));
|
||||
|
||||
REQUIRE(data == 1);
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipMemRangeGetAttribute_Positive_ReadMostly_Partial_Range") {
|
||||
if (!DeviceAttributesSupport(0, hipDeviceAttributeManagedMemory)) {
|
||||
HipTest::HIP_SKIP_TEST("Managed memory not supported");
|
||||
return;
|
||||
}
|
||||
|
||||
LinearAllocGuard<void> allocation(LinearAllocs::hipMallocManaged, 2 * kPageSize);
|
||||
|
||||
HIP_CHECK(hipMemAdvise(allocation.ptr(), kPageSize, hipMemAdviseSetReadMostly, 0));
|
||||
|
||||
int32_t data;
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(data), hipMemRangeAttributeReadMostly,
|
||||
allocation.ptr(), 2 * kPageSize));
|
||||
|
||||
REQUIRE(data == 0);
|
||||
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(data), hipMemRangeAttributeReadMostly,
|
||||
allocation.ptr(), kPageSize));
|
||||
|
||||
REQUIRE(data == 1);
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipMemRangeGetAttribute_Positive_PreferredLocation_Basic") {
|
||||
if (!DeviceAttributesSupport(0, hipDeviceAttributeManagedMemory)) {
|
||||
HipTest::HIP_SKIP_TEST("Managed memory not supported");
|
||||
return;
|
||||
}
|
||||
|
||||
LinearAllocGuard<void> allocation(LinearAllocs::hipMallocManaged, kPageSize);
|
||||
|
||||
int32_t data;
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(data), hipMemRangeAttributePreferredLocation,
|
||||
allocation.ptr(), kPageSize));
|
||||
|
||||
REQUIRE(data == hipInvalidDeviceId);
|
||||
|
||||
HIP_CHECK(hipMemAdvise(allocation.ptr(), kPageSize, hipMemAdviseSetPreferredLocation, 0));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(data), hipMemRangeAttributePreferredLocation,
|
||||
allocation.ptr(), kPageSize));
|
||||
|
||||
REQUIRE(data == 0);
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipMemRangeGetAttribute_Positive_PreferredLocation_CPU") {
|
||||
if (!DeviceAttributesSupport(0, hipDeviceAttributeManagedMemory)) {
|
||||
HipTest::HIP_SKIP_TEST("Managed memory not supported");
|
||||
return;
|
||||
}
|
||||
|
||||
LinearAllocGuard<void> allocation(LinearAllocs::hipMallocManaged, kPageSize);
|
||||
|
||||
HIP_CHECK(
|
||||
hipMemAdvise(allocation.ptr(), kPageSize, hipMemAdviseSetPreferredLocation, hipCpuDeviceId));
|
||||
|
||||
int32_t data;
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(data), hipMemRangeAttributePreferredLocation,
|
||||
allocation.ptr(), kPageSize));
|
||||
|
||||
REQUIRE(data == hipCpuDeviceId);
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipMemRangeGetAttribute_Positive_PreferredLocation_Partial_Range") {
|
||||
if (!DeviceAttributesSupport(0, hipDeviceAttributeManagedMemory)) {
|
||||
HipTest::HIP_SKIP_TEST("Managed memory not supported");
|
||||
return;
|
||||
}
|
||||
|
||||
LinearAllocGuard<void> allocation(LinearAllocs::hipMallocManaged, 2 * kPageSize);
|
||||
|
||||
HIP_CHECK(hipMemAdvise(allocation.ptr(), kPageSize, hipMemAdviseSetPreferredLocation, 0));
|
||||
|
||||
int32_t data;
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(data), hipMemRangeAttributePreferredLocation,
|
||||
allocation.ptr(), 2 * kPageSize));
|
||||
|
||||
REQUIRE(data == hipInvalidDeviceId);
|
||||
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(data), hipMemRangeAttributePreferredLocation,
|
||||
allocation.ptr(), kPageSize));
|
||||
|
||||
REQUIRE(data == 0);
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipMemRangeGetAttribute_Positive_LastPrefetchLocation_Basic") {
|
||||
if (!DeviceAttributesSupport(0, hipDeviceAttributeManagedMemory)) {
|
||||
HipTest::HIP_SKIP_TEST("Managed memory not supported");
|
||||
return;
|
||||
}
|
||||
|
||||
LinearAllocGuard<void> allocation(LinearAllocs::hipMallocManaged, kPageSize);
|
||||
|
||||
int32_t data;
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(data), hipMemRangeAttributeLastPrefetchLocation,
|
||||
allocation.ptr(), kPageSize));
|
||||
|
||||
REQUIRE(data == hipInvalidDeviceId);
|
||||
|
||||
HIP_CHECK(hipMemPrefetchAsync(allocation.ptr(), kPageSize, 0));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(data), hipMemRangeAttributeLastPrefetchLocation,
|
||||
allocation.ptr(), kPageSize));
|
||||
|
||||
REQUIRE(data == 0);
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipMemRangeGetAttribute_Positive_LastPrefetchLocation_CPU") {
|
||||
if (!DeviceAttributesSupport(0, hipDeviceAttributeManagedMemory)) {
|
||||
HipTest::HIP_SKIP_TEST("Managed memory not supported");
|
||||
return;
|
||||
}
|
||||
|
||||
LinearAllocGuard<void> allocation(LinearAllocs::hipMallocManaged, kPageSize);
|
||||
|
||||
HIP_CHECK(hipMemPrefetchAsync(allocation.ptr(), kPageSize, hipCpuDeviceId));
|
||||
|
||||
int32_t data;
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(data), hipMemRangeAttributeLastPrefetchLocation,
|
||||
allocation.ptr(), kPageSize));
|
||||
|
||||
REQUIRE(data == hipCpuDeviceId);
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipMemRangeGetAttribute_Positive_LastPrefetchLocation_Partial_Range") {
|
||||
if (!DeviceAttributesSupport(0, hipDeviceAttributeManagedMemory)) {
|
||||
HipTest::HIP_SKIP_TEST("Managed memory not supported");
|
||||
return;
|
||||
}
|
||||
|
||||
LinearAllocGuard<void> allocation(LinearAllocs::hipMallocManaged, 2 * kPageSize);
|
||||
|
||||
HIP_CHECK(hipMemPrefetchAsync(allocation.ptr(), kPageSize, 0));
|
||||
|
||||
int32_t data;
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(data), hipMemRangeAttributeLastPrefetchLocation,
|
||||
allocation.ptr(), 2 * kPageSize));
|
||||
|
||||
REQUIRE(data == hipInvalidDeviceId);
|
||||
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(data), hipMemRangeAttributeLastPrefetchLocation,
|
||||
allocation.ptr(), kPageSize));
|
||||
|
||||
REQUIRE(data == 0);
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipMemRangeGetAttribute_Positive_AccessedBy_Basic") {
|
||||
if (!DeviceAttributesSupport(0, hipDeviceAttributeManagedMemory)) {
|
||||
HipTest::HIP_SKIP_TEST("Managed memory not supported");
|
||||
return;
|
||||
}
|
||||
|
||||
LinearAllocGuard<void> allocation(LinearAllocs::hipMallocManaged, kPageSize);
|
||||
|
||||
std::array<int32_t, 4> data;
|
||||
HIP_CHECK(hipMemRangeGetAttribute(data.data(), sizeof(data), hipMemRangeAttributeAccessedBy,
|
||||
allocation.ptr(), kPageSize));
|
||||
|
||||
for (auto device : data) {
|
||||
REQUIRE(device == hipInvalidDeviceId);
|
||||
}
|
||||
|
||||
HIP_CHECK(hipMemAdvise(allocation.ptr(), kPageSize, hipMemAdviseSetAccessedBy, hipCpuDeviceId));
|
||||
HIP_CHECK(hipMemAdvise(allocation.ptr(), kPageSize, hipMemAdviseSetAccessedBy, 0));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(data.data(), sizeof(data), hipMemRangeAttributeAccessedBy,
|
||||
allocation.ptr(), kPageSize));
|
||||
|
||||
// Use std::find since there is no guaranteed order in which devices will be returned
|
||||
REQUIRE(std::find(cbegin(data), cend(data), hipCpuDeviceId) != cend(data));
|
||||
REQUIRE(std::find(cbegin(data), cend(data), 0) != cend(data));
|
||||
|
||||
// All the unused slots should be at the end
|
||||
for (auto it = cbegin(data) + 2; it != cend(data); ++it) {
|
||||
REQUIRE(*it == hipInvalidDeviceId);
|
||||
}
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipMemRangeGetAttribute_Positive_AccessedBy_Partial_Range") {
|
||||
if (!DeviceAttributesSupport(0, hipDeviceAttributeManagedMemory)) {
|
||||
HipTest::HIP_SKIP_TEST("Managed memory not supported");
|
||||
return;
|
||||
}
|
||||
|
||||
static int HmmAttrPrint() {
|
||||
int managed = 0;
|
||||
WARN("The following are the attribute values related to HMM for"
|
||||
" device 0:\n");
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed,
|
||||
hipDeviceAttributeDirectManagedMemAccessFromHost, 0));
|
||||
WARN("hipDeviceAttributeDirectManagedMemAccessFromHost: " << managed);
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed,
|
||||
hipDeviceAttributeConcurrentManagedAccess, 0));
|
||||
WARN("hipDeviceAttributeConcurrentManagedAccess: " << managed);
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed,
|
||||
hipDeviceAttributePageableMemoryAccess, 0));
|
||||
WARN("hipDeviceAttributePageableMemoryAccess: " << managed);
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed,
|
||||
hipDeviceAttributePageableMemoryAccessUsesHostPageTables, 0));
|
||||
WARN("hipDeviceAttributePageableMemoryAccessUsesHostPageTables:"
|
||||
<< managed);
|
||||
LinearAllocGuard<void> allocation(LinearAllocs::hipMallocManaged, 2 * kPageSize);
|
||||
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
|
||||
0));
|
||||
WARN("hipDeviceAttributeManagedMemory: " << managed);
|
||||
return managed;
|
||||
}
|
||||
HIP_CHECK(hipMemAdvise(allocation.ptr(), kPageSize, hipMemAdviseSetAccessedBy, hipCpuDeviceId));
|
||||
HIP_CHECK(hipMemAdvise(allocation.ptr(), kPageSize, hipMemAdviseSetAccessedBy, 0));
|
||||
|
||||
// The following function tests the count parameter(last param) to
|
||||
// hipMemRangeGetAttribute api by passing possible extreme values.
|
||||
// Curently the only way to test if count param working properly is to verify
|
||||
// the first parameter of hipMemRangeGetAttribute() api has value 1 stored
|
||||
TEST_CASE("Unit_hipMemRangeGetAttribute_TstCountParam") {
|
||||
int MangdMem = HmmAttrPrint();
|
||||
if (MangdMem == 1) {
|
||||
int MEM_SIZE = 4096, RND_NUM = 9999, FLG_READMOSTLY_ENBLD = 1;
|
||||
bool IfTestPassed = true;
|
||||
int data = RND_NUM, *devPtr = nullptr;
|
||||
size_t TotGpuMem, TotGpuFreeMem;
|
||||
HIP_CHECK(hipMemGetInfo(&TotGpuFreeMem, &TotGpuMem));
|
||||
std::array<int32_t, 4> data;
|
||||
HIP_CHECK(hipMemRangeGetAttribute(data.data(), sizeof(data), hipMemRangeAttributeAccessedBy,
|
||||
allocation.ptr(), 2 * kPageSize));
|
||||
|
||||
HIP_CHECK(hipMallocManaged(&devPtr, MEM_SIZE, hipMemAttachGlobal));
|
||||
HIP_CHECK(hipMemAdvise(devPtr, MEM_SIZE, hipMemAdviseSetReadMostly, 0));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(reinterpret_cast<void*>(&data),
|
||||
sizeof(int),
|
||||
hipMemRangeAttributeReadMostly,
|
||||
devPtr, MEM_SIZE));
|
||||
if (data != FLG_READMOSTLY_ENBLD) {
|
||||
WARN("hipMemRangeGetAttribute() api didnt return expected value!\n");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
HIP_CHECK(hipFree(devPtr));
|
||||
HIP_CHECK(hipMallocManaged(&devPtr, TotGpuFreeMem, hipMemAttachGlobal));
|
||||
HIP_CHECK(hipMemAdvise(devPtr, TotGpuFreeMem, hipMemAdviseSetReadMostly,
|
||||
0));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int),
|
||||
hipMemRangeAttributeReadMostly,
|
||||
devPtr, TotGpuFreeMem));
|
||||
for (auto device : data) {
|
||||
REQUIRE(device == hipInvalidDeviceId);
|
||||
}
|
||||
|
||||
if (data != FLG_READMOSTLY_ENBLD) {
|
||||
WARN("hipMemRangeGetAttribute() api didnt return expected value!\n");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
HIP_CHECK(hipFree(devPtr));
|
||||
HIP_CHECK(hipMallocManaged(&devPtr, (TotGpuFreeMem - 1),
|
||||
hipMemAttachGlobal));
|
||||
HIP_CHECK(hipMemAdvise(devPtr, (TotGpuFreeMem - 1),
|
||||
hipMemAdviseSetReadMostly, 0));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int),
|
||||
hipMemRangeAttributeReadMostly,
|
||||
devPtr, (TotGpuFreeMem - 1)));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(data.data(), sizeof(data), hipMemRangeAttributeAccessedBy,
|
||||
allocation.ptr(), kPageSize));
|
||||
|
||||
if (data != FLG_READMOSTLY_ENBLD) {
|
||||
WARN("hipMemRangeGetAttribute() api didnt return expected value!\n");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
HIP_CHECK(hipFree(devPtr));
|
||||
// Use std::find since there is no guaranteed order in which devices will be returned
|
||||
REQUIRE(std::find(cbegin(data), cend(data), hipCpuDeviceId) != cend(data));
|
||||
REQUIRE(std::find(cbegin(data), cend(data), 0) != cend(data));
|
||||
|
||||
REQUIRE(IfTestPassed);
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
// All the unused slots should be at the end
|
||||
for (auto it = cbegin(data) + 2; it != cend(data); ++it) {
|
||||
REQUIRE(*it == hipInvalidDeviceId);
|
||||
}
|
||||
}
|
||||
|
||||
/* The following Test Case does negative tests on hipMemRangeGetAttribute()*/
|
||||
TEST_CASE("Unit_hipMemRangeGetAttribute_Positive_AccessedBy_MultiDevice") {
|
||||
if (!DeviceAttributesSupport(0, hipDeviceAttributeManagedMemory)) {
|
||||
HipTest::HIP_SKIP_TEST("Managed memory not supported");
|
||||
return;
|
||||
}
|
||||
|
||||
TEST_CASE("Unit_hipMemRangeGetAttribute_NegativeTests") {
|
||||
int MangdMem = HmmAttrPrint();
|
||||
if (MangdMem == 1) {
|
||||
int MEM_SIZE = 4096, RND_NUM = 9999;
|
||||
float *devPtr = nullptr;
|
||||
int NumDevs;
|
||||
HIP_CHECK(hipGetDeviceCount(&NumDevs));
|
||||
int data = RND_NUM;
|
||||
int *OutData = new int[NumDevs];
|
||||
for (int m = 0; m < NumDevs; ++m) {
|
||||
OutData[m] = RND_NUM;
|
||||
}
|
||||
HIP_CHECK(hipMallocManaged(&devPtr, MEM_SIZE, hipMemAttachGlobal));
|
||||
HIP_CHECK(hipMemAdvise(devPtr, MEM_SIZE, hipMemAdviseSetReadMostly, 0));
|
||||
const auto device_count = HipTest::getDeviceCount();
|
||||
if (device_count < 2) {
|
||||
HipTest::HIP_SKIP_TEST("Two or more device are required");
|
||||
return;
|
||||
}
|
||||
|
||||
// checking the behavior with dataSize 0
|
||||
SECTION("checking the behavior with dataSize 0") {
|
||||
REQUIRE(CheckError(hipMemRangeGetAttribute(&data, 0,
|
||||
hipMemRangeAttributeReadMostly,
|
||||
devPtr, MEM_SIZE), __LINE__));
|
||||
}
|
||||
// checking the behavior with dataSize > 4 and even
|
||||
SECTION("checking the behavior with dataSize > 4 and even") {
|
||||
REQUIRE(CheckError(hipMemRangeGetAttribute(OutData, 6,
|
||||
hipMemRangeAttributeReadMostly,
|
||||
devPtr, MEM_SIZE), __LINE__));
|
||||
}
|
||||
// checking the behavior with dataSize > 4 and odd
|
||||
SECTION("checking the behavior with dataSize > 4 and odd") {
|
||||
REQUIRE(CheckError(hipMemRangeGetAttribute(OutData, 7,
|
||||
hipMemRangeAttributeReadMostly,
|
||||
devPtr, MEM_SIZE), __LINE__));
|
||||
}
|
||||
// checking the behavior with dataSize which is not multiple of 4
|
||||
SECTION("checking the behavior with dataSize which is not multiple of 4") {
|
||||
REQUIRE(CheckError(hipMemRangeGetAttribute(OutData, 27,
|
||||
hipMemRangeAttributeReadMostly,
|
||||
devPtr, MEM_SIZE), __LINE__));
|
||||
}
|
||||
// checking the behaviour with devPtr(4th param) as NULL
|
||||
SECTION("checking the behaviour with devPtr(4th param) as NULL") {
|
||||
REQUIRE(CheckError(hipMemRangeGetAttribute(&data, sizeof(int),
|
||||
hipMemRangeAttributeReadMostly,
|
||||
NULL, MEM_SIZE), __LINE__));
|
||||
}
|
||||
// checking the behaviour with count(5th param) as 0
|
||||
SECTION("checking the behaviour with count(5th param) as 0") {
|
||||
REQUIRE(CheckError(hipMemRangeGetAttribute(&data, sizeof(int),
|
||||
hipMemRangeAttributeReadMostly,
|
||||
devPtr, 0), __LINE__));
|
||||
}
|
||||
// checking the behavior with invalid attribute (3rd param) as 0
|
||||
// as it is attribute hence avoiding the negative tests with 3rd param
|
||||
LinearAllocGuard<void> allocation(LinearAllocs::hipMallocManaged, kPageSize);
|
||||
|
||||
// checking the behaviour of the api with ptr allocated using
|
||||
// hipHostMalloc
|
||||
void *ptr = nullptr;
|
||||
SECTION("Checking behavior with hipHostMalloc ptr") {
|
||||
HIP_CHECK(hipHostMalloc(&ptr, MEM_SIZE, 0));
|
||||
REQUIRE(CheckError(hipMemRangeGetAttribute(&data, sizeof(int),
|
||||
hipMemRangeAttributeReadMostly,
|
||||
ptr, MEM_SIZE), __LINE__));
|
||||
HIP_CHECK(hipHostFree(ptr));
|
||||
}
|
||||
HIP_CHECK(hipFree(devPtr));
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
std::vector<int32_t> data(device_count);
|
||||
HIP_CHECK(hipMemRangeGetAttribute(data.data(), sizeof(data), hipMemRangeAttributeAccessedBy,
|
||||
allocation.ptr(), kPageSize));
|
||||
|
||||
for (auto device : data) {
|
||||
REQUIRE(device == hipInvalidDeviceId);
|
||||
}
|
||||
|
||||
for (auto device = 0; device < device_count; ++device) {
|
||||
HIP_CHECK(hipMemAdvise(allocation.ptr(), kPageSize, hipMemAdviseSetAccessedBy, device));
|
||||
}
|
||||
|
||||
HIP_CHECK(hipMemRangeGetAttribute(data.data(), sizeof(data), hipMemRangeAttributeAccessedBy,
|
||||
allocation.ptr(), kPageSize));
|
||||
|
||||
// Use std::find since there is no guaranteed order in which devices will be returned
|
||||
for (auto device = 0; device < device_count; ++device) {
|
||||
REQUIRE(std::find(cbegin(data), cend(data), device) != cend(data));
|
||||
}
|
||||
}
|
||||
|
||||
/* This test case checks the behavior of hipMemRangeGetAttribute() with
|
||||
AccessedBy flag is consistent with cuda's counter part*/
|
||||
TEST_CASE("Unit_hipMemRangeGetAttribute_AccessedBy1") {
|
||||
int managed = HmmAttrPrint();
|
||||
if (managed == 1) {
|
||||
int Ngpus = 0, *Hmm = NULL, MEM_SZ = 4096, RND_NUM = 999;
|
||||
HIP_CHECK(hipGetDeviceCount(&Ngpus));
|
||||
int *OutData = new int[Ngpus];
|
||||
for (int i = 0; i < Ngpus; ++i) {
|
||||
OutData[Ngpus] = RND_NUM;
|
||||
}
|
||||
HIP_CHECK(hipMallocManaged(&Hmm, MEM_SZ));
|
||||
HIP_CHECK(hipMemAdvise(Hmm, MEM_SZ, hipMemAdviseSetAccessedBy, 0));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(OutData, 4*Ngpus,
|
||||
hipMemRangeAttributeAccessedBy,
|
||||
Hmm, MEM_SZ));
|
||||
if (OutData[0] != 0) {
|
||||
WARN("Didn't receive expected value at line: " << __LINE__);
|
||||
REQUIRE(false);
|
||||
}
|
||||
for (int i = 1; i < Ngpus; ++i) {
|
||||
if (OutData[i] != -2) {
|
||||
WARN("Didn't receive expected value at line: " << __LINE__);
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
if (Ngpus >= 2) {
|
||||
for (int i = 0; i < Ngpus; ++i) {
|
||||
HIP_CHECK(hipMemAdvise(Hmm, MEM_SZ, hipMemAdviseSetAccessedBy, i));
|
||||
}
|
||||
// checking the behavior with dataSize less than the number of gpus
|
||||
// This should not result in segfault.
|
||||
HIP_CHECK(hipMemRangeGetAttribute(OutData, 4*(Ngpus-1),
|
||||
hipMemRangeAttributeAccessedBy,
|
||||
Hmm, MEM_SZ));
|
||||
// OutData should have stored the gpu ordinals for which AccessedBy is
|
||||
// assigned except for the last element which should have -2 stored
|
||||
// so as to be consistent with cuda's behavior
|
||||
for (int i = 0; i < (Ngpus - 1); ++i) {
|
||||
if (OutData[i] != i) {
|
||||
WARN("Didn't receive expected value at line: " << __LINE__);
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
if (OutData[Ngpus - 1] != -2) {
|
||||
WARN("Didn't receive expected value at line: " << __LINE__);
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
HIP_CHECK(hipFree(Hmm));
|
||||
delete[] OutData;
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
TEST_CASE("Unit_hipMemRangeGetAttribute_Negative_Parameters") {
|
||||
if (!DeviceAttributesSupport(0, hipDeviceAttributeManagedMemory)) {
|
||||
HipTest::HIP_SKIP_TEST("Managed memory not supported");
|
||||
return;
|
||||
}
|
||||
|
||||
int32_t data;
|
||||
LinearAllocGuard<void> managed(LinearAllocs::hipMallocManaged, kPageSize);
|
||||
|
||||
SECTION("data == nullptr") {
|
||||
HIP_CHECK_ERROR(hipMemRangeGetAttribute(nullptr, 4, hipMemRangeAttributeReadMostly,
|
||||
managed.ptr(), kPageSize),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("data_size == 0") {
|
||||
HIP_CHECK_ERROR(
|
||||
hipMemRangeGetAttribute(&data, 0, hipMemRangeAttributeReadMostly, managed.ptr(), kPageSize),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("data_size != 4 with hipMemRangeAttributeReadMostly") {
|
||||
HIP_CHECK_ERROR(
|
||||
hipMemRangeGetAttribute(&data, 8, hipMemRangeAttributeReadMostly, managed.ptr(), kPageSize),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("data_size != 4 with hipMemRangeAttributePreferredLocation") {
|
||||
HIP_CHECK_ERROR(hipMemRangeGetAttribute(&data, 8, hipMemRangeAttributePreferredLocation,
|
||||
managed.ptr(), kPageSize),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("data_size != 4 with hipMemRangeAttributeLastPrefetchLocation") {
|
||||
HIP_CHECK_ERROR(hipMemRangeGetAttribute(&data, 8, hipMemRangeAttributeLastPrefetchLocation,
|
||||
managed.ptr(), kPageSize),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("data_size is not a multiple of 4 with hipMemRangeAttributeAccessedBy") {
|
||||
HIP_CHECK_ERROR(hipMemRangeGetAttribute(&data, 10, hipMemRangeAttributeAccessedBy,
|
||||
managed.ptr(), kPageSize),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("invalid attribute") {
|
||||
HIP_CHECK_ERROR(hipMemRangeGetAttribute(&data, 4, static_cast<hipMemRangeAttribute>(999),
|
||||
managed.ptr(), kPageSize),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("dev_ptr == nullptr") {
|
||||
HIP_CHECK_ERROR(
|
||||
hipMemRangeGetAttribute(&data, 4, hipMemRangeAttributeReadMostly, nullptr, kPageSize),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("dev_ptr is not managed memory") {
|
||||
LinearAllocGuard<void> non_managed(LinearAllocs::hipMalloc, kPageSize);
|
||||
HIP_CHECK_ERROR(hipMemRangeGetAttribute(&data, 4, hipMemRangeAttributeReadMostly,
|
||||
non_managed.ptr(), kPageSize),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("count == 0") {
|
||||
HIP_CHECK_ERROR(
|
||||
hipMemRangeGetAttribute(&data, 4, hipMemRangeAttributeReadMostly, managed.ptr(), 0),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
|
||||
|
||||
/* Allocate 4 * page size of memory with the flag hipMemAttachGloal. Advise
|
||||
AccessedBy, ReadMostly and PreferredLocation to first half(2*pageSz) of the
|
||||
memory and probe the for the flags which are set earlier using
|
||||
hipMemRangeGetAttribute() api for the full size(4*PageSz).*/
|
||||
/* Need to discuss the difference in behavior w.r.t cuda*/
|
||||
|
||||
TEST_CASE("Unit_hipMemRangeGetAttribte_3") {
|
||||
int managed = HmmAttrPrint();
|
||||
if (managed == 1) {
|
||||
int Ngpus = 0, *Hmm = NULL, MEM_SZ = 4096*4, RND_NUM = 999;
|
||||
HIP_CHECK(hipGetDeviceCount(&Ngpus));
|
||||
int *OutData = new int[Ngpus];
|
||||
for (int i = 0; i < Ngpus; ++i) {
|
||||
OutData[Ngpus] = RND_NUM;
|
||||
}
|
||||
HIP_CHECK(hipMallocManaged(&Hmm, MEM_SZ));
|
||||
HIP_CHECK(hipMemAdvise(Hmm, MEM_SZ/2, hipMemAdviseSetAccessedBy, 0));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(OutData, 4*Ngpus,
|
||||
hipMemRangeAttributeAccessedBy,
|
||||
(Hmm), MEM_SZ));
|
||||
|
||||
HIP_CHECK(hipMemAdvise(Hmm, MEM_SZ/2, hipMemAdviseSetReadMostly, 0));
|
||||
// The Api called below should not fail
|
||||
HIP_CHECK(hipMemRangeGetAttribute(OutData, 4,
|
||||
hipMemRangeAttributeReadMostly,
|
||||
(Hmm), MEM_SZ));
|
||||
|
||||
HIP_CHECK(hipMemAdvise(Hmm, MEM_SZ/2, hipMemAdviseSetPreferredLocation, 0));
|
||||
// The api called below should not fail
|
||||
HIP_CHECK(hipMemRangeGetAttribute(OutData, 4,
|
||||
hipMemRangeAttributePreferredLocation,
|
||||
(Hmm), MEM_SZ));
|
||||
HIP_CHECK(hipFree(Hmm));
|
||||
delete[] OutData;
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
/* The following scenarios tests that probing the attributes which are not set
|
||||
by hipMemAdvise() but being probed using hipMemRangeGetAttribute() should
|
||||
not result in a crash*/
|
||||
|
||||
TEST_CASE("Unit_hipMemRangeGetAttribute_4") {
|
||||
int managed = HmmAttrPrint();
|
||||
if (managed == 1) {
|
||||
int *Hmm = NULL, PageSz = 4096, Ngpus, RND_NUM = 999;
|
||||
HIP_CHECK(hipGetDeviceCount(&Ngpus));
|
||||
int *OutData = new int[Ngpus];
|
||||
for (int i = 0; i < Ngpus; ++i) {
|
||||
OutData[i] = RND_NUM;
|
||||
}
|
||||
HIP_CHECK(hipMallocManaged(&Hmm, 4*PageSz));
|
||||
SECTION("Set ReadMostly & probe other flags") {
|
||||
HIP_CHECK(hipMemAdvise(Hmm, 4*PageSz, hipMemAdviseSetReadMostly, 0));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(OutData, 4*Ngpus,
|
||||
hipMemRangeAttributeAccessedBy,
|
||||
Hmm, 4*PageSz));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(OutData, 4,
|
||||
hipMemRangeAttributePreferredLocation,
|
||||
Hmm, 4*PageSz));
|
||||
HIP_CHECK(hipMemAdvise(Hmm, 4*PageSz, hipMemAdviseUnsetReadMostly, 0));
|
||||
}
|
||||
SECTION("Set AccessedBy & probe other flags") {
|
||||
HIP_CHECK(hipMemAdvise(Hmm, 4*PageSz, hipMemAdviseSetAccessedBy, 0));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(OutData, 4,
|
||||
hipMemRangeAttributeReadMostly,
|
||||
Hmm, 4*PageSz));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(OutData, 4,
|
||||
hipMemRangeAttributePreferredLocation,
|
||||
Hmm, 4*PageSz));
|
||||
HIP_CHECK(hipMemAdvise(Hmm, 4*PageSz, hipMemAdviseUnsetAccessedBy, 0));
|
||||
}
|
||||
SECTION("Set AccessedBy & probe other flags") {
|
||||
HIP_CHECK(hipMemAdvise(Hmm, 4*PageSz, hipMemAdviseSetPreferredLocation,
|
||||
0));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(OutData, 4,
|
||||
hipMemRangeAttributeReadMostly,
|
||||
Hmm, 4*PageSz));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(OutData, 4*Ngpus,
|
||||
hipMemRangeAttributeAccessedBy,
|
||||
Hmm, 4*PageSz));
|
||||
HIP_CHECK(hipMemAdvise(Hmm, 4*PageSz, hipMemAdviseUnsetPreferredLocation,
|
||||
0));
|
||||
}
|
||||
HIP_CHECK(hipFree(Hmm));
|
||||
delete[] OutData;
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
/* The following scenario is a simple test which does the following:
|
||||
Allocate Hmm memory --> hipMemPrefetchAsync() to device 0 and then
|
||||
probe LastPrefetchLocation attribute using hipMemRangeGetAttribute*/
|
||||
|
||||
TEST_CASE("Unit_hipMemRangeGetAttribute_PrefetchAndGtAttr") {
|
||||
int managed = HmmAttrPrint();
|
||||
if (managed == 1) {
|
||||
int Ngpus = 0, *Hmm = NULL, RND_NUM = 999;
|
||||
size_t PageSz = 4096;
|
||||
HIP_CHECK(hipGetDeviceCount(&Ngpus));
|
||||
|
||||
int *OutData = new int[Ngpus];
|
||||
for (int i = 0; i < Ngpus; ++i) {
|
||||
OutData[Ngpus] = RND_NUM;
|
||||
}
|
||||
HIP_CHECK(hipMallocManaged(&Hmm, PageSz*4));
|
||||
hipStream_t strm;
|
||||
HIP_CHECK(hipStreamCreate(&strm));
|
||||
HIP_CHECK(hipMemPrefetchAsync(Hmm, PageSz*4, 0, strm));
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(OutData, 4,
|
||||
hipMemRangeAttributeLastPrefetchLocation,
|
||||
Hmm, PageSz*4));
|
||||
HIP_CHECK(hipStreamDestroy(strm));
|
||||
HIP_CHECK(hipFree(Hmm));
|
||||
if (OutData[0] != 0) {
|
||||
WARN("Didnt receive expected value at line: " << __LINE__);
|
||||
delete[] OutData;
|
||||
REQUIRE(false);
|
||||
}
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
@@ -0,0 +1,408 @@
|
||||
/*
|
||||
Copyright (c) 2021 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, INCLUDING 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 ANY 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.
|
||||
*/
|
||||
|
||||
/* Test Case Description:
|
||||
Scenario-1: The following function tests the count parameter(last param) to
|
||||
hipMemRangeGetAttribute api by passing possible extreme values.
|
||||
Curently the only way to test if count param working properly is to verify
|
||||
the first parameter of hipMemRangeGetAttribute() api has value 1 stored
|
||||
|
||||
Scenario-2: This test case checks the behavior of hipMemRangeGetAttribute() with
|
||||
AccessedBy flag is consistent with cuda's counter part
|
||||
|
||||
Scenario-3: Allocate 4 * page size of memory with the flag hipMemAttachGloal. Advise
|
||||
AccessedBy, ReadMostly and PreferredLocation to first half(2*pageSz) of the
|
||||
memory and probe the for the flags which are set earlier using
|
||||
hipMemRangeGetAttribute() api for the full size(4*PageSz).
|
||||
|
||||
|
||||
Scenario-4: The following scenarios tests that probing the attributes which are not set
|
||||
by hipMemAdvise() but being probed using hipMemRangeGetAttribute() should
|
||||
not result in a crash
|
||||
|
||||
Scenario-5: The following scenario is a simple test which does the following:
|
||||
Allocate Hmm memory --> hipMemPrefetchAsync() to device 0 and then
|
||||
probe LastPrefetchLocation attribute using hipMemRangeGetAttribute
|
||||
|
||||
Scenario-6: The following Test Case does negative tests on hipMemRangeGetAttribute()*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#include <stdlib.h>
|
||||
#ifdef __linux__
|
||||
#include <unistd.h>
|
||||
#include <sys/sysinfo.h>
|
||||
#endif
|
||||
|
||||
static bool CheckError(hipError_t err, int LineNo) {
|
||||
if (err == hipSuccess) {
|
||||
WARN("Error expected but received hipSuccess at line no.:"
|
||||
<< LineNo);
|
||||
return false;
|
||||
} else {
|
||||
return true;
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
static int HmmAttrPrint() {
|
||||
int managed = 0;
|
||||
WARN("The following are the attribute values related to HMM for"
|
||||
" device 0:\n");
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed,
|
||||
hipDeviceAttributeDirectManagedMemAccessFromHost, 0));
|
||||
WARN("hipDeviceAttributeDirectManagedMemAccessFromHost: " << managed);
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed,
|
||||
hipDeviceAttributeConcurrentManagedAccess, 0));
|
||||
WARN("hipDeviceAttributeConcurrentManagedAccess: " << managed);
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed,
|
||||
hipDeviceAttributePageableMemoryAccess, 0));
|
||||
WARN("hipDeviceAttributePageableMemoryAccess: " << managed);
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed,
|
||||
hipDeviceAttributePageableMemoryAccessUsesHostPageTables, 0));
|
||||
WARN("hipDeviceAttributePageableMemoryAccessUsesHostPageTables:"
|
||||
<< managed);
|
||||
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
|
||||
0));
|
||||
WARN("hipDeviceAttributeManagedMemory: " << managed);
|
||||
return managed;
|
||||
}
|
||||
|
||||
// The following function tests the count parameter(last param) to
|
||||
// hipMemRangeGetAttribute api by passing possible extreme values.
|
||||
// Curently the only way to test if count param working properly is to verify
|
||||
// the first parameter of hipMemRangeGetAttribute() api has value 1 stored
|
||||
TEST_CASE("Unit_hipMemRangeGetAttribute_TstCountParam") {
|
||||
int MangdMem = HmmAttrPrint();
|
||||
if (MangdMem == 1) {
|
||||
int MEM_SIZE = 4096, RND_NUM = 9999, FLG_READMOSTLY_ENBLD = 1;
|
||||
bool IfTestPassed = true;
|
||||
int data = RND_NUM, *devPtr = nullptr;
|
||||
size_t TotGpuMem, TotGpuFreeMem;
|
||||
HIP_CHECK(hipMemGetInfo(&TotGpuFreeMem, &TotGpuMem));
|
||||
|
||||
HIP_CHECK(hipMallocManaged(&devPtr, MEM_SIZE, hipMemAttachGlobal));
|
||||
HIP_CHECK(hipMemAdvise(devPtr, MEM_SIZE, hipMemAdviseSetReadMostly, 0));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(reinterpret_cast<void*>(&data),
|
||||
sizeof(int),
|
||||
hipMemRangeAttributeReadMostly,
|
||||
devPtr, MEM_SIZE));
|
||||
if (data != FLG_READMOSTLY_ENBLD) {
|
||||
WARN("hipMemRangeGetAttribute() api didnt return expected value!\n");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
HIP_CHECK(hipFree(devPtr));
|
||||
HIP_CHECK(hipMallocManaged(&devPtr, TotGpuFreeMem, hipMemAttachGlobal));
|
||||
HIP_CHECK(hipMemAdvise(devPtr, TotGpuFreeMem, hipMemAdviseSetReadMostly,
|
||||
0));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int),
|
||||
hipMemRangeAttributeReadMostly,
|
||||
devPtr, TotGpuFreeMem));
|
||||
|
||||
if (data != FLG_READMOSTLY_ENBLD) {
|
||||
WARN("hipMemRangeGetAttribute() api didnt return expected value!\n");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
HIP_CHECK(hipFree(devPtr));
|
||||
HIP_CHECK(hipMallocManaged(&devPtr, (TotGpuFreeMem - 1),
|
||||
hipMemAttachGlobal));
|
||||
HIP_CHECK(hipMemAdvise(devPtr, (TotGpuFreeMem - 1),
|
||||
hipMemAdviseSetReadMostly, 0));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int),
|
||||
hipMemRangeAttributeReadMostly,
|
||||
devPtr, (TotGpuFreeMem - 1)));
|
||||
|
||||
if (data != FLG_READMOSTLY_ENBLD) {
|
||||
WARN("hipMemRangeGetAttribute() api didnt return expected value!\n");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
HIP_CHECK(hipFree(devPtr));
|
||||
|
||||
REQUIRE(IfTestPassed);
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
}
|
||||
|
||||
/* The following Test Case does negative tests on hipMemRangeGetAttribute()*/
|
||||
|
||||
TEST_CASE("Unit_hipMemRangeGetAttribute_NegativeTests") {
|
||||
int MangdMem = HmmAttrPrint();
|
||||
if (MangdMem == 1) {
|
||||
int MEM_SIZE = 4096, RND_NUM = 9999;
|
||||
float *devPtr = nullptr;
|
||||
int NumDevs;
|
||||
HIP_CHECK(hipGetDeviceCount(&NumDevs));
|
||||
int data = RND_NUM;
|
||||
int *OutData = new int[NumDevs];
|
||||
for (int m = 0; m < NumDevs; ++m) {
|
||||
OutData[m] = RND_NUM;
|
||||
}
|
||||
HIP_CHECK(hipMallocManaged(&devPtr, MEM_SIZE, hipMemAttachGlobal));
|
||||
HIP_CHECK(hipMemAdvise(devPtr, MEM_SIZE, hipMemAdviseSetReadMostly, 0));
|
||||
|
||||
// checking the behavior with dataSize 0
|
||||
SECTION("checking the behavior with dataSize 0") {
|
||||
REQUIRE(CheckError(hipMemRangeGetAttribute(&data, 0,
|
||||
hipMemRangeAttributeReadMostly,
|
||||
devPtr, MEM_SIZE), __LINE__));
|
||||
}
|
||||
// checking the behavior with dataSize > 4 and even
|
||||
SECTION("checking the behavior with dataSize > 4 and even") {
|
||||
REQUIRE(CheckError(hipMemRangeGetAttribute(OutData, 6,
|
||||
hipMemRangeAttributeReadMostly,
|
||||
devPtr, MEM_SIZE), __LINE__));
|
||||
}
|
||||
// checking the behavior with dataSize > 4 and odd
|
||||
SECTION("checking the behavior with dataSize > 4 and odd") {
|
||||
REQUIRE(CheckError(hipMemRangeGetAttribute(OutData, 7,
|
||||
hipMemRangeAttributeReadMostly,
|
||||
devPtr, MEM_SIZE), __LINE__));
|
||||
}
|
||||
// checking the behavior with dataSize which is not multiple of 4
|
||||
SECTION("checking the behavior with dataSize which is not multiple of 4") {
|
||||
REQUIRE(CheckError(hipMemRangeGetAttribute(OutData, 27,
|
||||
hipMemRangeAttributeReadMostly,
|
||||
devPtr, MEM_SIZE), __LINE__));
|
||||
}
|
||||
// checking the behaviour with devPtr(4th param) as NULL
|
||||
SECTION("checking the behaviour with devPtr(4th param) as NULL") {
|
||||
REQUIRE(CheckError(hipMemRangeGetAttribute(&data, sizeof(int),
|
||||
hipMemRangeAttributeReadMostly,
|
||||
NULL, MEM_SIZE), __LINE__));
|
||||
}
|
||||
// checking the behaviour with count(5th param) as 0
|
||||
SECTION("checking the behaviour with count(5th param) as 0") {
|
||||
REQUIRE(CheckError(hipMemRangeGetAttribute(&data, sizeof(int),
|
||||
hipMemRangeAttributeReadMostly,
|
||||
devPtr, 0), __LINE__));
|
||||
}
|
||||
// checking the behavior with invalid attribute (3rd param) as 0
|
||||
// as it is attribute hence avoiding the negative tests with 3rd param
|
||||
|
||||
// checking the behaviour of the api with ptr allocated using
|
||||
// hipHostMalloc
|
||||
void *ptr = nullptr;
|
||||
SECTION("Checking behavior with hipHostMalloc ptr") {
|
||||
HIP_CHECK(hipHostMalloc(&ptr, MEM_SIZE, 0));
|
||||
REQUIRE(CheckError(hipMemRangeGetAttribute(&data, sizeof(int),
|
||||
hipMemRangeAttributeReadMostly,
|
||||
ptr, MEM_SIZE), __LINE__));
|
||||
HIP_CHECK(hipHostFree(ptr));
|
||||
}
|
||||
HIP_CHECK(hipFree(devPtr));
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
}
|
||||
|
||||
/* This test case checks the behavior of hipMemRangeGetAttribute() with
|
||||
AccessedBy flag is consistent with cuda's counter part*/
|
||||
TEST_CASE("Unit_hipMemRangeGetAttribute_AccessedBy1") {
|
||||
int managed = HmmAttrPrint();
|
||||
if (managed == 1) {
|
||||
int Ngpus = 0, *Hmm = NULL, MEM_SZ = 4096, RND_NUM = 999;
|
||||
HIP_CHECK(hipGetDeviceCount(&Ngpus));
|
||||
int *OutData = new int[Ngpus];
|
||||
for (int i = 0; i < Ngpus; ++i) {
|
||||
OutData[Ngpus] = RND_NUM;
|
||||
}
|
||||
HIP_CHECK(hipMallocManaged(&Hmm, MEM_SZ));
|
||||
HIP_CHECK(hipMemAdvise(Hmm, MEM_SZ, hipMemAdviseSetAccessedBy, 0));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(OutData, 4*Ngpus,
|
||||
hipMemRangeAttributeAccessedBy,
|
||||
Hmm, MEM_SZ));
|
||||
if (OutData[0] != 0) {
|
||||
WARN("Didn't receive expected value at line: " << __LINE__);
|
||||
REQUIRE(false);
|
||||
}
|
||||
for (int i = 1; i < Ngpus; ++i) {
|
||||
if (OutData[i] != -2) {
|
||||
WARN("Didn't receive expected value at line: " << __LINE__);
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
if (Ngpus >= 2) {
|
||||
for (int i = 0; i < Ngpus; ++i) {
|
||||
HIP_CHECK(hipMemAdvise(Hmm, MEM_SZ, hipMemAdviseSetAccessedBy, i));
|
||||
}
|
||||
// checking the behavior with dataSize less than the number of gpus
|
||||
// This should not result in segfault.
|
||||
HIP_CHECK(hipMemRangeGetAttribute(OutData, 4*(Ngpus-1),
|
||||
hipMemRangeAttributeAccessedBy,
|
||||
Hmm, MEM_SZ));
|
||||
// OutData should have stored the gpu ordinals for which AccessedBy is
|
||||
// assigned except for the last element which should have -2 stored
|
||||
// so as to be consistent with cuda's behavior
|
||||
for (int i = 0; i < (Ngpus - 1); ++i) {
|
||||
if (OutData[i] != i) {
|
||||
WARN("Didn't receive expected value at line: " << __LINE__);
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
if (OutData[Ngpus - 1] != -2) {
|
||||
WARN("Didn't receive expected value at line: " << __LINE__);
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
HIP_CHECK(hipFree(Hmm));
|
||||
delete[] OutData;
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
|
||||
|
||||
/* Allocate 4 * page size of memory with the flag hipMemAttachGloal. Advise
|
||||
AccessedBy, ReadMostly and PreferredLocation to first half(2*pageSz) of the
|
||||
memory and probe the for the flags which are set earlier using
|
||||
hipMemRangeGetAttribute() api for the full size(4*PageSz).*/
|
||||
/* Need to discuss the difference in behavior w.r.t cuda*/
|
||||
|
||||
TEST_CASE("Unit_hipMemRangeGetAttribte_3") {
|
||||
int managed = HmmAttrPrint();
|
||||
if (managed == 1) {
|
||||
int Ngpus = 0, *Hmm = NULL, MEM_SZ = 4096*4, RND_NUM = 999;
|
||||
HIP_CHECK(hipGetDeviceCount(&Ngpus));
|
||||
int *OutData = new int[Ngpus];
|
||||
for (int i = 0; i < Ngpus; ++i) {
|
||||
OutData[Ngpus] = RND_NUM;
|
||||
}
|
||||
HIP_CHECK(hipMallocManaged(&Hmm, MEM_SZ));
|
||||
HIP_CHECK(hipMemAdvise(Hmm, MEM_SZ/2, hipMemAdviseSetAccessedBy, 0));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(OutData, 4*Ngpus,
|
||||
hipMemRangeAttributeAccessedBy,
|
||||
(Hmm), MEM_SZ));
|
||||
|
||||
HIP_CHECK(hipMemAdvise(Hmm, MEM_SZ/2, hipMemAdviseSetReadMostly, 0));
|
||||
// The Api called below should not fail
|
||||
HIP_CHECK(hipMemRangeGetAttribute(OutData, 4,
|
||||
hipMemRangeAttributeReadMostly,
|
||||
(Hmm), MEM_SZ));
|
||||
|
||||
HIP_CHECK(hipMemAdvise(Hmm, MEM_SZ/2, hipMemAdviseSetPreferredLocation, 0));
|
||||
// The api called below should not fail
|
||||
HIP_CHECK(hipMemRangeGetAttribute(OutData, 4,
|
||||
hipMemRangeAttributePreferredLocation,
|
||||
(Hmm), MEM_SZ));
|
||||
HIP_CHECK(hipFree(Hmm));
|
||||
delete[] OutData;
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
/* The following scenarios tests that probing the attributes which are not set
|
||||
by hipMemAdvise() but being probed using hipMemRangeGetAttribute() should
|
||||
not result in a crash*/
|
||||
|
||||
TEST_CASE("Unit_hipMemRangeGetAttribute_4") {
|
||||
int managed = HmmAttrPrint();
|
||||
if (managed == 1) {
|
||||
int *Hmm = NULL, PageSz = 4096, Ngpus, RND_NUM = 999;
|
||||
HIP_CHECK(hipGetDeviceCount(&Ngpus));
|
||||
int *OutData = new int[Ngpus];
|
||||
for (int i = 0; i < Ngpus; ++i) {
|
||||
OutData[i] = RND_NUM;
|
||||
}
|
||||
HIP_CHECK(hipMallocManaged(&Hmm, 4*PageSz));
|
||||
SECTION("Set ReadMostly & probe other flags") {
|
||||
HIP_CHECK(hipMemAdvise(Hmm, 4*PageSz, hipMemAdviseSetReadMostly, 0));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(OutData, 4*Ngpus,
|
||||
hipMemRangeAttributeAccessedBy,
|
||||
Hmm, 4*PageSz));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(OutData, 4,
|
||||
hipMemRangeAttributePreferredLocation,
|
||||
Hmm, 4*PageSz));
|
||||
HIP_CHECK(hipMemAdvise(Hmm, 4*PageSz, hipMemAdviseUnsetReadMostly, 0));
|
||||
}
|
||||
SECTION("Set AccessedBy & probe other flags") {
|
||||
HIP_CHECK(hipMemAdvise(Hmm, 4*PageSz, hipMemAdviseSetAccessedBy, 0));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(OutData, 4,
|
||||
hipMemRangeAttributeReadMostly,
|
||||
Hmm, 4*PageSz));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(OutData, 4,
|
||||
hipMemRangeAttributePreferredLocation,
|
||||
Hmm, 4*PageSz));
|
||||
HIP_CHECK(hipMemAdvise(Hmm, 4*PageSz, hipMemAdviseUnsetAccessedBy, 0));
|
||||
}
|
||||
SECTION("Set AccessedBy & probe other flags") {
|
||||
HIP_CHECK(hipMemAdvise(Hmm, 4*PageSz, hipMemAdviseSetPreferredLocation,
|
||||
0));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(OutData, 4,
|
||||
hipMemRangeAttributeReadMostly,
|
||||
Hmm, 4*PageSz));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(OutData, 4*Ngpus,
|
||||
hipMemRangeAttributeAccessedBy,
|
||||
Hmm, 4*PageSz));
|
||||
HIP_CHECK(hipMemAdvise(Hmm, 4*PageSz, hipMemAdviseUnsetPreferredLocation,
|
||||
0));
|
||||
}
|
||||
HIP_CHECK(hipFree(Hmm));
|
||||
delete[] OutData;
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
/* The following scenario is a simple test which does the following:
|
||||
Allocate Hmm memory --> hipMemPrefetchAsync() to device 0 and then
|
||||
probe LastPrefetchLocation attribute using hipMemRangeGetAttribute*/
|
||||
|
||||
TEST_CASE("Unit_hipMemRangeGetAttribute_PrefetchAndGtAttr") {
|
||||
int managed = HmmAttrPrint();
|
||||
if (managed == 1) {
|
||||
int Ngpus = 0, *Hmm = NULL, RND_NUM = 999;
|
||||
size_t PageSz = 4096;
|
||||
HIP_CHECK(hipGetDeviceCount(&Ngpus));
|
||||
|
||||
int *OutData = new int[Ngpus];
|
||||
for (int i = 0; i < Ngpus; ++i) {
|
||||
OutData[Ngpus] = RND_NUM;
|
||||
}
|
||||
HIP_CHECK(hipMallocManaged(&Hmm, PageSz*4));
|
||||
hipStream_t strm;
|
||||
HIP_CHECK(hipStreamCreate(&strm));
|
||||
HIP_CHECK(hipMemPrefetchAsync(Hmm, PageSz*4, 0, strm));
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
HIP_CHECK(hipMemRangeGetAttribute(OutData, 4,
|
||||
hipMemRangeAttributeLastPrefetchLocation,
|
||||
Hmm, PageSz*4));
|
||||
HIP_CHECK(hipStreamDestroy(strm));
|
||||
HIP_CHECK(hipFree(Hmm));
|
||||
if (OutData[0] != 0) {
|
||||
WARN("Didnt receive expected value at line: " << __LINE__);
|
||||
delete[] OutData;
|
||||
REQUIRE(false);
|
||||
}
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
}
|
||||
|
||||
@@ -1,325 +1,155 @@
|
||||
/*
|
||||
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 WARRANNTY OF ANY KIND, EXPRESS OR
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
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 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
|
||||
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.
|
||||
*/
|
||||
|
||||
/* Test Case Description:
|
||||
Scenario-1: Testing basic working of hipMemRangeGetAttributes()
|
||||
api with different flags
|
||||
Scenario-2: Negative testing with hipMemRangeGetAttributes() api
|
||||
*/
|
||||
|
||||
#include <hip/hip_runtime_api.h>
|
||||
#include <hip_test_common.hh>
|
||||
#define MEM_SIZE 8192
|
||||
#include <resource_guards.hh>
|
||||
#include <utils.hh>
|
||||
|
||||
static bool CheckError(hipError_t err, int LineNo) {
|
||||
if (err == hipSuccess) {
|
||||
WARN("Error expected but received hipSuccess at line no.:"
|
||||
<< LineNo);
|
||||
return false;
|
||||
} else {
|
||||
return true;
|
||||
TEST_CASE("Unit_hipMemRangeGetAttributes_Positive_Basic") {
|
||||
if (!DeviceAttributesSupport(0, hipDeviceAttributeManagedMemory)) {
|
||||
HipTest::HIP_SKIP_TEST("Managed memory not supported");
|
||||
return;
|
||||
}
|
||||
|
||||
LinearAllocGuard<void> allocation(LinearAllocs::hipMallocManaged, kPageSize);
|
||||
|
||||
HIP_CHECK(hipMemAdvise(allocation.ptr(), kPageSize, hipMemAdviseSetReadMostly, 0));
|
||||
HIP_CHECK(hipMemAdvise(allocation.ptr(), kPageSize, hipMemAdviseSetPreferredLocation, 0));
|
||||
HIP_CHECK(hipMemPrefetchAsync(allocation.ptr(), kPageSize, hipCpuDeviceId));
|
||||
HIP_CHECK(hipMemAdvise(allocation.ptr(), kPageSize, hipMemAdviseSetAccessedBy, 0));
|
||||
|
||||
constexpr size_t num_attributes = 4;
|
||||
std::array<hipMemRangeAttribute, num_attributes> attributes = {
|
||||
hipMemRangeAttributeReadMostly, hipMemRangeAttributePreferredLocation,
|
||||
hipMemRangeAttributeLastPrefetchLocation, hipMemRangeAttributeAccessedBy};
|
||||
|
||||
std::array<int32_t*, num_attributes> data;
|
||||
for (auto& ptr : data) {
|
||||
ptr = new int32_t;
|
||||
}
|
||||
std::array<size_t, num_attributes> data_sizes = {4, 4, 4, 4};
|
||||
|
||||
HIP_CHECK(hipMemRangeGetAttributes(reinterpret_cast<void**>(data.data()), data_sizes.data(),
|
||||
attributes.data(), num_attributes, allocation.ptr(),
|
||||
kPageSize));
|
||||
|
||||
REQUIRE(data[0][0] == 1);
|
||||
REQUIRE(data[1][0] == 0);
|
||||
REQUIRE(data[2][0] == hipCpuDeviceId);
|
||||
REQUIRE(data[3][0] == 0);
|
||||
|
||||
for (auto ptr : data) {
|
||||
delete ptr;
|
||||
}
|
||||
}
|
||||
|
||||
static int HmmAttrPrint() {
|
||||
int managed = 0;
|
||||
WARN("The following are the attribute values related to HMM for"
|
||||
" device 0:\n");
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed,
|
||||
hipDeviceAttributeDirectManagedMemAccessFromHost, 0));
|
||||
WARN("hipDeviceAttributeDirectManagedMemAccessFromHost: " << managed);
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed,
|
||||
hipDeviceAttributeConcurrentManagedAccess, 0));
|
||||
WARN("hipDeviceAttributeConcurrentManagedAccess: " << managed);
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed,
|
||||
hipDeviceAttributePageableMemoryAccess, 0));
|
||||
WARN("hipDeviceAttributePageableMemoryAccess: " << managed);
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed,
|
||||
hipDeviceAttributePageableMemoryAccessUsesHostPageTables, 0));
|
||||
WARN("hipDeviceAttributePageableMemoryAccessUsesHostPageTables:"
|
||||
<< managed);
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
|
||||
0));
|
||||
WARN("hipDeviceAttributeManagedMemory: " << managed);
|
||||
return managed;
|
||||
}
|
||||
TEST_CASE("Unit_hipMemRangeGetAttributes_Negative_Parameters") {
|
||||
if (!DeviceAttributesSupport(0, hipDeviceAttributeManagedMemory)) {
|
||||
HipTest::HIP_SKIP_TEST("Managed memory not supported");
|
||||
return;
|
||||
}
|
||||
|
||||
#ifdef __linux__
|
||||
/* Test Scenario: Testing basic working of hipMemRangeGetAttributes()
|
||||
api with different flags */
|
||||
constexpr size_t num_attributes = 4;
|
||||
hipMemRangeAttribute attributes[] = {
|
||||
hipMemRangeAttributeReadMostly, hipMemRangeAttributePreferredLocation,
|
||||
hipMemRangeAttributeLastPrefetchLocation, hipMemRangeAttributeAccessedBy};
|
||||
|
||||
TEST_CASE("Unit_hipMemRangeGetAttributes_TstFlgs") {
|
||||
int MangdMem = HmmAttrPrint();
|
||||
if (MangdMem == 1) {
|
||||
bool IfTestPassed = true;
|
||||
int NumDevs = 0;
|
||||
int *Outpt[4], *AcsdBy = nullptr;
|
||||
float *Hmm = nullptr;
|
||||
hipStream_t strm;
|
||||
hipMemRangeAttribute AttrArr[4] =
|
||||
{hipMemRangeAttributeReadMostly,
|
||||
hipMemRangeAttributePreferredLocation,
|
||||
hipMemRangeAttributeAccessedBy,
|
||||
hipMemRangeAttributeLastPrefetchLocation};
|
||||
HIP_CHECK(hipGetDeviceCount(&NumDevs));
|
||||
AcsdBy = new int(NumDevs);
|
||||
size_t dataSizes[4] = {sizeof(int), sizeof(int),
|
||||
(NumDevs * sizeof(int)), sizeof(int)};
|
||||
Outpt[0] = new int;
|
||||
Outpt[1] = new int;
|
||||
Outpt[2] = new int[NumDevs];
|
||||
Outpt[3] = new int;
|
||||
HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE, hipMemAttachGlobal));
|
||||
for (int i = 0; i < NumDevs; ++i) {
|
||||
HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetReadMostly, i));
|
||||
HIP_CHECK(hipMemRangeGetAttributes(reinterpret_cast<void**>(Outpt),
|
||||
dataSizes, AttrArr, 4, Hmm,
|
||||
MEM_SIZE));
|
||||
if (*(Outpt[0]) != 1) {
|
||||
WARN("Attempt to set hipMemAdviseSetReadMostly flag failed!\n");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseUnsetReadMostly, i));
|
||||
HIP_CHECK(hipMemRangeGetAttributes(reinterpret_cast<void**>(Outpt),
|
||||
reinterpret_cast<size_t*>(dataSizes),
|
||||
AttrArr, 4, Hmm, MEM_SIZE));
|
||||
int32_t* data[num_attributes];
|
||||
for (auto& ptr : data) {
|
||||
ptr = new int32_t;
|
||||
}
|
||||
size_t data_sizes[] = {4, 4, 4, 4};
|
||||
|
||||
if (*(Outpt[0]) != 0) {
|
||||
WARN("Attempt to set hipMemAdviseUnsetReadMostly flag failed!\n");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
LinearAllocGuard<void> managed(LinearAllocs::hipMallocManaged, kPageSize);
|
||||
|
||||
HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE,
|
||||
hipMemAdviseSetPreferredLocation, i));
|
||||
HIP_CHECK(hipMemRangeGetAttributes(reinterpret_cast<void**>(Outpt),
|
||||
reinterpret_cast<size_t*>(dataSizes),
|
||||
AttrArr, 4, Hmm, MEM_SIZE));
|
||||
if (*(Outpt[1]) != i) {
|
||||
WARN("Attempt to set hipMemAdviseSetPreferredLocation flag");
|
||||
WARN(" failed!\n");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetAccessedBy, i));
|
||||
HIP_CHECK(hipMemRangeGetAttributes(reinterpret_cast<void**>(Outpt),
|
||||
reinterpret_cast<size_t*>(dataSizes),
|
||||
AttrArr, 4, Hmm, MEM_SIZE));
|
||||
if ((Outpt[2][0]) != i) {
|
||||
WARN("Attempt to set hipMemAdviseSetAccessedBy flag");
|
||||
WARN(" failed!\n");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
SECTION("data == nullptr") {
|
||||
HIP_CHECK_ERROR(hipMemRangeGetAttributes(nullptr, data_sizes, attributes, num_attributes,
|
||||
managed.ptr(), kPageSize),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseUnsetAccessedBy, i));
|
||||
HIP_CHECK(hipMemRangeGetAttributes(reinterpret_cast<void**>(Outpt),
|
||||
reinterpret_cast<size_t*>(dataSizes),
|
||||
AttrArr, 4, Hmm, MEM_SIZE));
|
||||
if (!((Outpt[2][i]) < 0)) {
|
||||
WARN("Attempt to set hipMemAdviseUnsetAccessedBy flag failed!\n");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
HIP_CHECK(hipStreamCreate(&strm));
|
||||
HIP_CHECK(hipMemPrefetchAsync(Hmm, MEM_SIZE, i, strm));
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
HIP_CHECK(hipMemRangeGetAttributes(reinterpret_cast<void**>(Outpt),
|
||||
reinterpret_cast<size_t*>(dataSizes),
|
||||
AttrArr, 4, Hmm, MEM_SIZE));
|
||||
if (*(Outpt[3]) != i) {
|
||||
WARN("Attempt to prefetch memory to device: " << i);
|
||||
WARN("failed!\n");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
// Prefetching back to Host
|
||||
HIP_CHECK(hipMemPrefetchAsync(Hmm, MEM_SIZE, -1, strm));
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
HIP_CHECK(hipMemRangeGetAttributes(reinterpret_cast<void**>(Outpt),
|
||||
reinterpret_cast<size_t*>(dataSizes),
|
||||
AttrArr, 4, Hmm, MEM_SIZE));
|
||||
if (*(Outpt[3]) != -1) {
|
||||
WARN("Attempt to prefetch memory to Host failed!\n");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
}
|
||||
SECTION("data contains invalid pointers") {
|
||||
void* invalid_data[num_attributes] = {nullptr};
|
||||
HIP_CHECK_ERROR(hipMemRangeGetAttributes(invalid_data, data_sizes, attributes, num_attributes,
|
||||
managed.ptr(), kPageSize),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
HIP_CHECK(hipFree(Hmm));
|
||||
delete[] AcsdBy;
|
||||
for (int i = 0; i < 4; ++i) {
|
||||
delete Outpt[i];
|
||||
}
|
||||
REQUIRE(IfTestPassed);
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
SECTION("data_sizes == nullptr") {
|
||||
HIP_CHECK_ERROR(hipMemRangeGetAttributes(reinterpret_cast<void**>(data), nullptr, attributes,
|
||||
num_attributes, managed.ptr(), kPageSize),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("data_sizes contains invalid values") {
|
||||
size_t invalid_data_sizes[] = {4, 5, 4, 6};
|
||||
HIP_CHECK_ERROR(hipMemRangeGetAttributes(reinterpret_cast<void**>(data), invalid_data_sizes,
|
||||
attributes, num_attributes, managed.ptr(), kPageSize),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("attributes == nullptr") {
|
||||
HIP_CHECK_ERROR(hipMemRangeGetAttributes(reinterpret_cast<void**>(data), data_sizes, nullptr,
|
||||
num_attributes, managed.ptr(), kPageSize),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("attributes contains invalid attributes") {
|
||||
hipMemRangeAttribute invalid_attributes[] = {
|
||||
hipMemRangeAttributeReadMostly, hipMemRangeAttributePreferredLocation,
|
||||
static_cast<hipMemRangeAttribute>(999), hipMemRangeAttributeAccessedBy};
|
||||
HIP_CHECK_ERROR(
|
||||
hipMemRangeGetAttributes(reinterpret_cast<void**>(data), data_sizes, invalid_attributes,
|
||||
num_attributes, managed.ptr(), kPageSize),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("num_attributes == 0") {
|
||||
HIP_CHECK_ERROR(hipMemRangeGetAttributes(reinterpret_cast<void**>(data), data_sizes, attributes,
|
||||
0, managed.ptr(), kPageSize),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("dev_ptr == nullptr") {
|
||||
HIP_CHECK_ERROR(hipMemRangeGetAttributes(reinterpret_cast<void**>(data), data_sizes, attributes,
|
||||
num_attributes, nullptr, kPageSize),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("dev_ptr is not managed memory") {
|
||||
LinearAllocGuard<void> non_managed(LinearAllocs::hipMalloc, kPageSize);
|
||||
HIP_CHECK_ERROR(hipMemRangeGetAttributes(reinterpret_cast<void**>(data), data_sizes, attributes,
|
||||
num_attributes, non_managed.ptr(), kPageSize),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
SECTION("count == 0") {
|
||||
HIP_CHECK_ERROR(hipMemRangeGetAttributes(reinterpret_cast<void**>(data), data_sizes, attributes,
|
||||
num_attributes, managed.ptr(), 0),
|
||||
hipErrorInvalidValue);
|
||||
}
|
||||
|
||||
for (auto ptr : data) {
|
||||
delete ptr;
|
||||
}
|
||||
}
|
||||
|
||||
/* Test Scenario: Negative testing with hipMemRangeGetAttributes() api*/
|
||||
TEST_CASE("Unit_hipMemRangeGetAttributes_NegativeTst") {
|
||||
int MangdMem = HmmAttrPrint();
|
||||
if (MangdMem == 1) {
|
||||
bool IfTestPassed = true;
|
||||
int NumDevs = 0, *Outpt[4];
|
||||
float *Hmm = nullptr;
|
||||
hipMemRangeAttribute AttrArr[4] =
|
||||
{hipMemRangeAttributeReadMostly,
|
||||
hipMemRangeAttributePreferredLocation,
|
||||
hipMemRangeAttributeAccessedBy,
|
||||
hipMemRangeAttributeLastPrefetchLocation};
|
||||
HIP_CHECK(hipGetDeviceCount(&NumDevs));
|
||||
size_t dataSizes[4] = {sizeof(int), sizeof(int),
|
||||
(NumDevs * sizeof(int)), sizeof(int)};
|
||||
Outpt[0] = new int;
|
||||
Outpt[1] = new int;
|
||||
Outpt[2] = new int[NumDevs];
|
||||
Outpt[3] = new int;
|
||||
HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE, hipMemAttachGlobal));
|
||||
HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE, hipMemAdviseSetReadMostly, 0));
|
||||
// passing zero for num of attributes param(4th)
|
||||
SECTION("passing zero for num of attributes param(4th)") {
|
||||
if (!CheckError(hipMemRangeGetAttributes(
|
||||
reinterpret_cast<void**>(Outpt),
|
||||
reinterpret_cast<size_t*>(dataSizes),
|
||||
AttrArr, 0, Hmm, MEM_SIZE), __LINE__)) {
|
||||
IfTestPassed = false;
|
||||
}
|
||||
}
|
||||
|
||||
// the first dataSize element passed as 0
|
||||
dataSizes[0] = 0;
|
||||
dataSizes[1] = sizeof(int);
|
||||
dataSizes[2] = NumDevs * sizeof(int);
|
||||
dataSizes[3] = sizeof(int);
|
||||
SECTION("the first dataSize element passed as 0") {
|
||||
if (!CheckError(hipMemRangeGetAttributes(
|
||||
reinterpret_cast<void**>(Outpt),
|
||||
reinterpret_cast<size_t*>(dataSizes),
|
||||
AttrArr, 4, Hmm, MEM_SIZE),
|
||||
__LINE__)) {
|
||||
IfTestPassed = false;
|
||||
}
|
||||
}
|
||||
// passing datasize as 2 while the requirement is multiple of 4
|
||||
dataSizes[0] = 2;
|
||||
dataSizes[1] = sizeof(int);
|
||||
dataSizes[2] = NumDevs * sizeof(int);
|
||||
dataSizes[3] = sizeof(int);
|
||||
SECTION("datasize as 2 while the requirement is multiple of 4") {
|
||||
if (!CheckError(hipMemRangeGetAttributes(
|
||||
reinterpret_cast<void**>(Outpt),
|
||||
reinterpret_cast<size_t*>(dataSizes),
|
||||
AttrArr, 4, Hmm, MEM_SIZE),
|
||||
__LINE__)) {
|
||||
IfTestPassed = false;
|
||||
}
|
||||
}
|
||||
// passing datasize as 6 while the requirement is multiple of 4
|
||||
dataSizes[0] = 6;
|
||||
dataSizes[1] = sizeof(int);
|
||||
dataSizes[2] = NumDevs * sizeof(int);
|
||||
dataSizes[3] = sizeof(int);
|
||||
SECTION("datasize as 6 while the requirement is multiple of 4") {
|
||||
if (!CheckError(hipMemRangeGetAttributes(
|
||||
reinterpret_cast<void**>(Outpt),
|
||||
reinterpret_cast<size_t*>(dataSizes),
|
||||
AttrArr, 4, Hmm, MEM_SIZE),
|
||||
__LINE__)) {
|
||||
IfTestPassed = false;
|
||||
}
|
||||
}
|
||||
// passing datasize as 7 while the requirement is multiple of 4
|
||||
dataSizes[0] = 7;
|
||||
dataSizes[1] = sizeof(int);
|
||||
dataSizes[2] = NumDevs * sizeof(int);
|
||||
dataSizes[3] = sizeof(int);
|
||||
SECTION("datasize as 7 while the requirement is multiple of 4") {
|
||||
if (!CheckError(hipMemRangeGetAttributes(
|
||||
reinterpret_cast<void**>(Outpt),
|
||||
reinterpret_cast<size_t*>(dataSizes),
|
||||
AttrArr, 4, Hmm, MEM_SIZE),
|
||||
__LINE__)) {
|
||||
IfTestPassed = false;
|
||||
}
|
||||
}
|
||||
// passing dataSize as 7 for attribute hipMemRangeAttributeAccessedBy
|
||||
hipMemRangeAttribute AttrArr1[1] = {hipMemRangeAttributeAccessedBy};
|
||||
dataSizes[2] = {7};
|
||||
SECTION("passing dataSize as 7 for attribute hipMemRangeAttrAccessedBy") {
|
||||
if (!CheckError(hipMemRangeGetAttributes(
|
||||
reinterpret_cast<void**>(Outpt),
|
||||
reinterpret_cast<size_t*>(dataSizes),
|
||||
AttrArr1, 1, Hmm, MEM_SIZE), __LINE__)) {
|
||||
IfTestPassed = false;
|
||||
}
|
||||
}
|
||||
// Passing NULL as first parameter
|
||||
SECTION("Passing NULL as first parameter") {
|
||||
if (!CheckError(hipMemRangeGetAttributes(
|
||||
reinterpret_cast<void**>(NULL),
|
||||
reinterpret_cast<size_t*>(dataSizes),
|
||||
AttrArr, 4, Hmm, MEM_SIZE),
|
||||
__LINE__)) {
|
||||
IfTestPassed = false;
|
||||
}
|
||||
}
|
||||
// Passing count parameter as zero
|
||||
SECTION("Passing count parameter as zero") {
|
||||
if (!CheckError(hipMemRangeGetAttributes(
|
||||
reinterpret_cast<void**>(Outpt),
|
||||
reinterpret_cast<size_t*>(dataSizes),
|
||||
AttrArr, 4, Hmm, 0),
|
||||
__LINE__)) {
|
||||
IfTestPassed = false;
|
||||
}
|
||||
}
|
||||
// Passing NULL for Attribute array(3rd param)
|
||||
SECTION("Passing NULL for Attribute array(3rd param)") {
|
||||
if (!CheckError(hipMemRangeGetAttributes(
|
||||
reinterpret_cast<void**>(Outpt),
|
||||
reinterpret_cast<size_t*>(dataSizes),
|
||||
NULL, 4, Hmm, MEM_SIZE),
|
||||
__LINE__)) {
|
||||
IfTestPassed = false;
|
||||
}
|
||||
}
|
||||
// Passing 0 for Attribute array(3rd param)
|
||||
SECTION("Passing 0 for Attribute array(3rd param)") {
|
||||
if (!CheckError(hipMemRangeGetAttributes(
|
||||
reinterpret_cast<void**>(Outpt),
|
||||
reinterpret_cast<size_t*>(dataSizes),
|
||||
0, 4, Hmm, MEM_SIZE),
|
||||
__LINE__)) {
|
||||
IfTestPassed = false;
|
||||
}
|
||||
}
|
||||
for (int i = 0; i < 4; ++i) {
|
||||
delete Outpt[i];
|
||||
}
|
||||
REQUIRE(IfTestPassed);
|
||||
|
||||
// The following scenarios have been removed considering the nature of the
|
||||
// api. With Consultation with Maneesh Gupta, the following scenarios
|
||||
// have been removed.
|
||||
// passing numAttributes as 4 while the attributes array has only 2 members
|
||||
// passing numAttributes as 10 while the attributes array has only 2 members
|
||||
// length of the list of dataSizes less than the number of
|
||||
// attributes being probed
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
}
|
||||
#endif
|
||||
|
||||
@@ -0,0 +1,325 @@
|
||||
/*
|
||||
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 WARRANNTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING 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 ANY 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.
|
||||
*/
|
||||
|
||||
/* Test Case Description:
|
||||
Scenario-1: Testing basic working of hipMemRangeGetAttributes()
|
||||
api with different flags
|
||||
Scenario-2: Negative testing with hipMemRangeGetAttributes() api
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#define MEM_SIZE 8192
|
||||
|
||||
static bool CheckError(hipError_t err, int LineNo) {
|
||||
if (err == hipSuccess) {
|
||||
WARN("Error expected but received hipSuccess at line no.:"
|
||||
<< LineNo);
|
||||
return false;
|
||||
} else {
|
||||
return true;
|
||||
}
|
||||
}
|
||||
|
||||
static int HmmAttrPrint() {
|
||||
int managed = 0;
|
||||
WARN("The following are the attribute values related to HMM for"
|
||||
" device 0:\n");
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed,
|
||||
hipDeviceAttributeDirectManagedMemAccessFromHost, 0));
|
||||
WARN("hipDeviceAttributeDirectManagedMemAccessFromHost: " << managed);
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed,
|
||||
hipDeviceAttributeConcurrentManagedAccess, 0));
|
||||
WARN("hipDeviceAttributeConcurrentManagedAccess: " << managed);
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed,
|
||||
hipDeviceAttributePageableMemoryAccess, 0));
|
||||
WARN("hipDeviceAttributePageableMemoryAccess: " << managed);
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed,
|
||||
hipDeviceAttributePageableMemoryAccessUsesHostPageTables, 0));
|
||||
WARN("hipDeviceAttributePageableMemoryAccessUsesHostPageTables:"
|
||||
<< managed);
|
||||
HIP_CHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeManagedMemory,
|
||||
0));
|
||||
WARN("hipDeviceAttributeManagedMemory: " << managed);
|
||||
return managed;
|
||||
}
|
||||
|
||||
#ifdef __linux__
|
||||
/* Test Scenario: Testing basic working of hipMemRangeGetAttributes()
|
||||
api with different flags */
|
||||
|
||||
TEST_CASE("Unit_hipMemRangeGetAttributes_TstFlgs") {
|
||||
int MangdMem = HmmAttrPrint();
|
||||
if (MangdMem == 1) {
|
||||
bool IfTestPassed = true;
|
||||
int NumDevs = 0;
|
||||
int *Outpt[4], *AcsdBy = nullptr;
|
||||
float *Hmm = nullptr;
|
||||
hipStream_t strm;
|
||||
hipMemRangeAttribute AttrArr[4] =
|
||||
{hipMemRangeAttributeReadMostly,
|
||||
hipMemRangeAttributePreferredLocation,
|
||||
hipMemRangeAttributeAccessedBy,
|
||||
hipMemRangeAttributeLastPrefetchLocation};
|
||||
HIP_CHECK(hipGetDeviceCount(&NumDevs));
|
||||
AcsdBy = new int(NumDevs);
|
||||
size_t dataSizes[4] = {sizeof(int), sizeof(int),
|
||||
(NumDevs * sizeof(int)), sizeof(int)};
|
||||
Outpt[0] = new int;
|
||||
Outpt[1] = new int;
|
||||
Outpt[2] = new int[NumDevs];
|
||||
Outpt[3] = new int;
|
||||
HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE, hipMemAttachGlobal));
|
||||
for (int i = 0; i < NumDevs; ++i) {
|
||||
HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetReadMostly, i));
|
||||
HIP_CHECK(hipMemRangeGetAttributes(reinterpret_cast<void**>(Outpt),
|
||||
dataSizes, AttrArr, 4, Hmm,
|
||||
MEM_SIZE));
|
||||
if (*(Outpt[0]) != 1) {
|
||||
WARN("Attempt to set hipMemAdviseSetReadMostly flag failed!\n");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseUnsetReadMostly, i));
|
||||
HIP_CHECK(hipMemRangeGetAttributes(reinterpret_cast<void**>(Outpt),
|
||||
reinterpret_cast<size_t*>(dataSizes),
|
||||
AttrArr, 4, Hmm, MEM_SIZE));
|
||||
|
||||
if (*(Outpt[0]) != 0) {
|
||||
WARN("Attempt to set hipMemAdviseUnsetReadMostly flag failed!\n");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
|
||||
HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE,
|
||||
hipMemAdviseSetPreferredLocation, i));
|
||||
HIP_CHECK(hipMemRangeGetAttributes(reinterpret_cast<void**>(Outpt),
|
||||
reinterpret_cast<size_t*>(dataSizes),
|
||||
AttrArr, 4, Hmm, MEM_SIZE));
|
||||
if (*(Outpt[1]) != i) {
|
||||
WARN("Attempt to set hipMemAdviseSetPreferredLocation flag");
|
||||
WARN(" failed!\n");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetAccessedBy, i));
|
||||
HIP_CHECK(hipMemRangeGetAttributes(reinterpret_cast<void**>(Outpt),
|
||||
reinterpret_cast<size_t*>(dataSizes),
|
||||
AttrArr, 4, Hmm, MEM_SIZE));
|
||||
if ((Outpt[2][0]) != i) {
|
||||
WARN("Attempt to set hipMemAdviseSetAccessedBy flag");
|
||||
WARN(" failed!\n");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
|
||||
HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseUnsetAccessedBy, i));
|
||||
HIP_CHECK(hipMemRangeGetAttributes(reinterpret_cast<void**>(Outpt),
|
||||
reinterpret_cast<size_t*>(dataSizes),
|
||||
AttrArr, 4, Hmm, MEM_SIZE));
|
||||
if (!((Outpt[2][i]) < 0)) {
|
||||
WARN("Attempt to set hipMemAdviseUnsetAccessedBy flag failed!\n");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
HIP_CHECK(hipStreamCreate(&strm));
|
||||
HIP_CHECK(hipMemPrefetchAsync(Hmm, MEM_SIZE, i, strm));
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
HIP_CHECK(hipMemRangeGetAttributes(reinterpret_cast<void**>(Outpt),
|
||||
reinterpret_cast<size_t*>(dataSizes),
|
||||
AttrArr, 4, Hmm, MEM_SIZE));
|
||||
if (*(Outpt[3]) != i) {
|
||||
WARN("Attempt to prefetch memory to device: " << i);
|
||||
WARN("failed!\n");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
// Prefetching back to Host
|
||||
HIP_CHECK(hipMemPrefetchAsync(Hmm, MEM_SIZE, -1, strm));
|
||||
HIP_CHECK(hipStreamSynchronize(strm));
|
||||
HIP_CHECK(hipMemRangeGetAttributes(reinterpret_cast<void**>(Outpt),
|
||||
reinterpret_cast<size_t*>(dataSizes),
|
||||
AttrArr, 4, Hmm, MEM_SIZE));
|
||||
if (*(Outpt[3]) != -1) {
|
||||
WARN("Attempt to prefetch memory to Host failed!\n");
|
||||
IfTestPassed = false;
|
||||
}
|
||||
}
|
||||
|
||||
HIP_CHECK(hipFree(Hmm));
|
||||
delete[] AcsdBy;
|
||||
for (int i = 0; i < 4; ++i) {
|
||||
delete Outpt[i];
|
||||
}
|
||||
REQUIRE(IfTestPassed);
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
}
|
||||
|
||||
/* Test Scenario: Negative testing with hipMemRangeGetAttributes() api*/
|
||||
TEST_CASE("Unit_hipMemRangeGetAttributes_NegativeTst") {
|
||||
int MangdMem = HmmAttrPrint();
|
||||
if (MangdMem == 1) {
|
||||
bool IfTestPassed = true;
|
||||
int NumDevs = 0, *Outpt[4];
|
||||
float *Hmm = nullptr;
|
||||
hipMemRangeAttribute AttrArr[4] =
|
||||
{hipMemRangeAttributeReadMostly,
|
||||
hipMemRangeAttributePreferredLocation,
|
||||
hipMemRangeAttributeAccessedBy,
|
||||
hipMemRangeAttributeLastPrefetchLocation};
|
||||
HIP_CHECK(hipGetDeviceCount(&NumDevs));
|
||||
size_t dataSizes[4] = {sizeof(int), sizeof(int),
|
||||
(NumDevs * sizeof(int)), sizeof(int)};
|
||||
Outpt[0] = new int;
|
||||
Outpt[1] = new int;
|
||||
Outpt[2] = new int[NumDevs];
|
||||
Outpt[3] = new int;
|
||||
HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE, hipMemAttachGlobal));
|
||||
HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE, hipMemAdviseSetReadMostly, 0));
|
||||
// passing zero for num of attributes param(4th)
|
||||
SECTION("passing zero for num of attributes param(4th)") {
|
||||
if (!CheckError(hipMemRangeGetAttributes(
|
||||
reinterpret_cast<void**>(Outpt),
|
||||
reinterpret_cast<size_t*>(dataSizes),
|
||||
AttrArr, 0, Hmm, MEM_SIZE), __LINE__)) {
|
||||
IfTestPassed = false;
|
||||
}
|
||||
}
|
||||
|
||||
// the first dataSize element passed as 0
|
||||
dataSizes[0] = 0;
|
||||
dataSizes[1] = sizeof(int);
|
||||
dataSizes[2] = NumDevs * sizeof(int);
|
||||
dataSizes[3] = sizeof(int);
|
||||
SECTION("the first dataSize element passed as 0") {
|
||||
if (!CheckError(hipMemRangeGetAttributes(
|
||||
reinterpret_cast<void**>(Outpt),
|
||||
reinterpret_cast<size_t*>(dataSizes),
|
||||
AttrArr, 4, Hmm, MEM_SIZE),
|
||||
__LINE__)) {
|
||||
IfTestPassed = false;
|
||||
}
|
||||
}
|
||||
// passing datasize as 2 while the requirement is multiple of 4
|
||||
dataSizes[0] = 2;
|
||||
dataSizes[1] = sizeof(int);
|
||||
dataSizes[2] = NumDevs * sizeof(int);
|
||||
dataSizes[3] = sizeof(int);
|
||||
SECTION("datasize as 2 while the requirement is multiple of 4") {
|
||||
if (!CheckError(hipMemRangeGetAttributes(
|
||||
reinterpret_cast<void**>(Outpt),
|
||||
reinterpret_cast<size_t*>(dataSizes),
|
||||
AttrArr, 4, Hmm, MEM_SIZE),
|
||||
__LINE__)) {
|
||||
IfTestPassed = false;
|
||||
}
|
||||
}
|
||||
// passing datasize as 6 while the requirement is multiple of 4
|
||||
dataSizes[0] = 6;
|
||||
dataSizes[1] = sizeof(int);
|
||||
dataSizes[2] = NumDevs * sizeof(int);
|
||||
dataSizes[3] = sizeof(int);
|
||||
SECTION("datasize as 6 while the requirement is multiple of 4") {
|
||||
if (!CheckError(hipMemRangeGetAttributes(
|
||||
reinterpret_cast<void**>(Outpt),
|
||||
reinterpret_cast<size_t*>(dataSizes),
|
||||
AttrArr, 4, Hmm, MEM_SIZE),
|
||||
__LINE__)) {
|
||||
IfTestPassed = false;
|
||||
}
|
||||
}
|
||||
// passing datasize as 7 while the requirement is multiple of 4
|
||||
dataSizes[0] = 7;
|
||||
dataSizes[1] = sizeof(int);
|
||||
dataSizes[2] = NumDevs * sizeof(int);
|
||||
dataSizes[3] = sizeof(int);
|
||||
SECTION("datasize as 7 while the requirement is multiple of 4") {
|
||||
if (!CheckError(hipMemRangeGetAttributes(
|
||||
reinterpret_cast<void**>(Outpt),
|
||||
reinterpret_cast<size_t*>(dataSizes),
|
||||
AttrArr, 4, Hmm, MEM_SIZE),
|
||||
__LINE__)) {
|
||||
IfTestPassed = false;
|
||||
}
|
||||
}
|
||||
// passing dataSize as 7 for attribute hipMemRangeAttributeAccessedBy
|
||||
hipMemRangeAttribute AttrArr1[1] = {hipMemRangeAttributeAccessedBy};
|
||||
dataSizes[2] = {7};
|
||||
SECTION("passing dataSize as 7 for attribute hipMemRangeAttrAccessedBy") {
|
||||
if (!CheckError(hipMemRangeGetAttributes(
|
||||
reinterpret_cast<void**>(Outpt),
|
||||
reinterpret_cast<size_t*>(dataSizes),
|
||||
AttrArr1, 1, Hmm, MEM_SIZE), __LINE__)) {
|
||||
IfTestPassed = false;
|
||||
}
|
||||
}
|
||||
// Passing NULL as first parameter
|
||||
SECTION("Passing NULL as first parameter") {
|
||||
if (!CheckError(hipMemRangeGetAttributes(
|
||||
reinterpret_cast<void**>(NULL),
|
||||
reinterpret_cast<size_t*>(dataSizes),
|
||||
AttrArr, 4, Hmm, MEM_SIZE),
|
||||
__LINE__)) {
|
||||
IfTestPassed = false;
|
||||
}
|
||||
}
|
||||
// Passing count parameter as zero
|
||||
SECTION("Passing count parameter as zero") {
|
||||
if (!CheckError(hipMemRangeGetAttributes(
|
||||
reinterpret_cast<void**>(Outpt),
|
||||
reinterpret_cast<size_t*>(dataSizes),
|
||||
AttrArr, 4, Hmm, 0),
|
||||
__LINE__)) {
|
||||
IfTestPassed = false;
|
||||
}
|
||||
}
|
||||
// Passing NULL for Attribute array(3rd param)
|
||||
SECTION("Passing NULL for Attribute array(3rd param)") {
|
||||
if (!CheckError(hipMemRangeGetAttributes(
|
||||
reinterpret_cast<void**>(Outpt),
|
||||
reinterpret_cast<size_t*>(dataSizes),
|
||||
NULL, 4, Hmm, MEM_SIZE),
|
||||
__LINE__)) {
|
||||
IfTestPassed = false;
|
||||
}
|
||||
}
|
||||
// Passing 0 for Attribute array(3rd param)
|
||||
SECTION("Passing 0 for Attribute array(3rd param)") {
|
||||
if (!CheckError(hipMemRangeGetAttributes(
|
||||
reinterpret_cast<void**>(Outpt),
|
||||
reinterpret_cast<size_t*>(dataSizes),
|
||||
0, 4, Hmm, MEM_SIZE),
|
||||
__LINE__)) {
|
||||
IfTestPassed = false;
|
||||
}
|
||||
}
|
||||
for (int i = 0; i < 4; ++i) {
|
||||
delete Outpt[i];
|
||||
}
|
||||
REQUIRE(IfTestPassed);
|
||||
|
||||
// The following scenarios have been removed considering the nature of the
|
||||
// api. With Consultation with Maneesh Gupta, the following scenarios
|
||||
// have been removed.
|
||||
// passing numAttributes as 4 while the attributes array has only 2 members
|
||||
// passing numAttributes as 10 while the attributes array has only 2 members
|
||||
// length of the list of dataSizes less than the number of
|
||||
// attributes being probed
|
||||
} else {
|
||||
SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
|
||||
"attribute. Hence skipping the testing with Pass result.\n");
|
||||
}
|
||||
}
|
||||
#endif
|
||||
Referencia en una nueva incidencia
Block a user