From 848418ec9b9769ecb15db5829f3564aa9ea8b102 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Mirza=20Halil=C4=8Devi=C4=87?= <109971222+mirza-halilcevic@users.noreply.github.com> Date: Tue, 17 Jan 2023 12:57:58 +0100 Subject: [PATCH] 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: 203b99423062610df25bfb5b075a4379f2424267] --- .../catch/include/resource_guards.hh | 57 +- .../catch/unit/memory/CMakeLists.txt | 4 + .../unit/memory/hipMemRangeGetAttribute.cc | 693 ++++++++---------- .../memory/hipMemRangeGetAttribute_old.cc | 408 +++++++++++ .../unit/memory/hipMemRangeGetAttributes.cc | 424 ++++------- .../memory/hipMemRangeGetAttributes_old.cc | 325 ++++++++ 6 files changed, 1241 insertions(+), 670 deletions(-) create mode 100644 projects/hip-tests/catch/unit/memory/hipMemRangeGetAttribute_old.cc create mode 100644 projects/hip-tests/catch/unit/memory/hipMemRangeGetAttributes_old.cc diff --git a/projects/hip-tests/catch/include/resource_guards.hh b/projects/hip-tests/catch/include/resource_guards.hh index d50ae05baa..af4c285e2a 100644 --- a/projects/hip-tests/catch/include/resource_guards.hh +++ b/projects/hip-tests/catch/include/resource_guards.hh @@ -196,11 +196,11 @@ template 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(hipEventDestroy(e)); + } + + hipEvent_t& operator[](int index) { return events_[index]; } + + operator hipEvent_t() const { return events_.at(0); } + + std::vector& event_list() { return events_; } + +private: + std::vector 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(hipStreamDestroy(s)); + } + + hipStream_t& operator[](int index) { return streams_[index]; } + + operator hipStream_t() const { return streams_.at(0); } + + std::vector& stream_list() { return streams_; } + +private: + std::vector streams_; +}; diff --git a/projects/hip-tests/catch/unit/memory/CMakeLists.txt b/projects/hip-tests/catch/unit/memory/CMakeLists.txt index 962f493a60..32ea04d3c1 100644 --- a/projects/hip-tests/catch/unit/memory/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/memory/CMakeLists.txt @@ -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 ) diff --git a/projects/hip-tests/catch/unit/memory/hipMemRangeGetAttribute.cc b/projects/hip-tests/catch/unit/memory/hipMemRangeGetAttribute.cc index 64fc8e6be4..447a3a0b56 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemRangeGetAttribute.cc +++ b/projects/hip-tests/catch/unit/memory/hipMemRangeGetAttribute.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 #include -#include -#ifdef __linux__ - #include - #include -#endif +#include +#include -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 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 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 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 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 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 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 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 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 allocation(LinearAllocs::hipMallocManaged, kPageSize); + + std::array 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 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 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(&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 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 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 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(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 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"); - } -} - diff --git a/projects/hip-tests/catch/unit/memory/hipMemRangeGetAttribute_old.cc b/projects/hip-tests/catch/unit/memory/hipMemRangeGetAttribute_old.cc new file mode 100644 index 0000000000..64fc8e6be4 --- /dev/null +++ b/projects/hip-tests/catch/unit/memory/hipMemRangeGetAttribute_old.cc @@ -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 +#include +#ifdef __linux__ + #include + #include +#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(&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"); + } +} + diff --git a/projects/hip-tests/catch/unit/memory/hipMemRangeGetAttributes.cc b/projects/hip-tests/catch/unit/memory/hipMemRangeGetAttributes.cc index e1240f4330..a3a9a5bf2a 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemRangeGetAttributes.cc +++ b/projects/hip-tests/catch/unit/memory/hipMemRangeGetAttributes.cc @@ -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 #include -#define MEM_SIZE 8192 +#include +#include -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 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 attributes = { + hipMemRangeAttributeReadMostly, hipMemRangeAttributePreferredLocation, + hipMemRangeAttributeLastPrefetchLocation, hipMemRangeAttributeAccessedBy}; + + std::array data; + for (auto& ptr : data) { + ptr = new int32_t; + } + std::array data_sizes = {4, 4, 4, 4}; + + HIP_CHECK(hipMemRangeGetAttributes(reinterpret_cast(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(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(Outpt), - reinterpret_cast(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 managed(LinearAllocs::hipMallocManaged, kPageSize); - HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, - hipMemAdviseSetPreferredLocation, i)); - HIP_CHECK(hipMemRangeGetAttributes(reinterpret_cast(Outpt), - reinterpret_cast(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(Outpt), - reinterpret_cast(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(Outpt), - reinterpret_cast(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(Outpt), - reinterpret_cast(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(Outpt), - reinterpret_cast(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(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(data), invalid_data_sizes, + attributes, num_attributes, managed.ptr(), kPageSize), + hipErrorInvalidValue); + } + + SECTION("attributes == nullptr") { + HIP_CHECK_ERROR(hipMemRangeGetAttributes(reinterpret_cast(data), data_sizes, nullptr, + num_attributes, managed.ptr(), kPageSize), + hipErrorInvalidValue); + } + + SECTION("attributes contains invalid attributes") { + hipMemRangeAttribute invalid_attributes[] = { + hipMemRangeAttributeReadMostly, hipMemRangeAttributePreferredLocation, + static_cast(999), hipMemRangeAttributeAccessedBy}; + HIP_CHECK_ERROR( + hipMemRangeGetAttributes(reinterpret_cast(data), data_sizes, invalid_attributes, + num_attributes, managed.ptr(), kPageSize), + hipErrorInvalidValue); + } + + SECTION("num_attributes == 0") { + HIP_CHECK_ERROR(hipMemRangeGetAttributes(reinterpret_cast(data), data_sizes, attributes, + 0, managed.ptr(), kPageSize), + hipErrorInvalidValue); + } + + SECTION("dev_ptr == nullptr") { + HIP_CHECK_ERROR(hipMemRangeGetAttributes(reinterpret_cast(data), data_sizes, attributes, + num_attributes, nullptr, kPageSize), + hipErrorInvalidValue); + } + + SECTION("dev_ptr is not managed memory") { + LinearAllocGuard non_managed(LinearAllocs::hipMalloc, kPageSize); + HIP_CHECK_ERROR(hipMemRangeGetAttributes(reinterpret_cast(data), data_sizes, attributes, + num_attributes, non_managed.ptr(), kPageSize), + hipErrorInvalidValue); + } + + SECTION("count == 0") { + HIP_CHECK_ERROR(hipMemRangeGetAttributes(reinterpret_cast(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(Outpt), - reinterpret_cast(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(Outpt), - reinterpret_cast(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(Outpt), - reinterpret_cast(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(Outpt), - reinterpret_cast(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(Outpt), - reinterpret_cast(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(Outpt), - reinterpret_cast(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(NULL), - reinterpret_cast(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(Outpt), - reinterpret_cast(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(Outpt), - reinterpret_cast(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(Outpt), - reinterpret_cast(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 diff --git a/projects/hip-tests/catch/unit/memory/hipMemRangeGetAttributes_old.cc b/projects/hip-tests/catch/unit/memory/hipMemRangeGetAttributes_old.cc new file mode 100644 index 0000000000..e1240f4330 --- /dev/null +++ b/projects/hip-tests/catch/unit/memory/hipMemRangeGetAttributes_old.cc @@ -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 +#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(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(Outpt), + reinterpret_cast(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(Outpt), + reinterpret_cast(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(Outpt), + reinterpret_cast(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(Outpt), + reinterpret_cast(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(Outpt), + reinterpret_cast(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(Outpt), + reinterpret_cast(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(Outpt), + reinterpret_cast(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(Outpt), + reinterpret_cast(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(Outpt), + reinterpret_cast(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(Outpt), + reinterpret_cast(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(Outpt), + reinterpret_cast(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(Outpt), + reinterpret_cast(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(NULL), + reinterpret_cast(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(Outpt), + reinterpret_cast(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(Outpt), + reinterpret_cast(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(Outpt), + reinterpret_cast(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