From 6b48c2f9195bd0e3aadf036a494e2584509679b6 Mon Sep 17 00:00:00 2001 From: agunashe <86270081+agunashe@users.noreply.github.com> Date: Wed, 1 Feb 2023 10:48:11 +0530 Subject: [PATCH 01/13] SWDEV-327563 - Disable Unit_hipGraphRetainUserObject_Functional_2 test. (#149) Intermittent issue: corrupted doubly linked list --- catch/hipTestMain/config/config_amd_linux_common.json | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/catch/hipTestMain/config/config_amd_linux_common.json b/catch/hipTestMain/config/config_amd_linux_common.json index b5eb2dfca6..be119a27ab 100644 --- a/catch/hipTestMain/config/config_amd_linux_common.json +++ b/catch/hipTestMain/config/config_amd_linux_common.json @@ -27,6 +27,8 @@ "Unit_hipStreamAttachMemAsync_Negative_Parameters", "Unit_hipMemGetAddressRange_Positive", "Unit_hipGraphAddMemcpyNode1D_Negative_Basic", - "Unit_hipStreamGetCaptureInfo_Nullstream_CaptureInfo" + "Unit_hipStreamGetCaptureInfo_Nullstream_CaptureInfo", + "intermittent issue: corrupted double-linked list", + "Unit_hipGraphRetainUserObject_Functional_2" ] } From cc0302dda31c229b6b3b645d91d0186056d6ddbe Mon Sep 17 00:00:00 2001 From: music-dino <111048524+music-dino@users.noreply.github.com> Date: Wed, 1 Feb 2023 07:20:14 +0100 Subject: [PATCH 02/13] EXSWHTEC-83 - Implement new and reimplement existing tests for hipMemAdvise (#16) - Basic positive tests - Negative parameter tests --- .../config/config_amd_linux_common.json | 2 + .../config/config_amd_windows_common.json | 2 + catch/unit/memory/CMakeLists.txt | 2 + catch/unit/memory/hipMemAdvise.cc | 1187 ++++------------- catch/unit/memory/hipMemAdvise_old.cc | 941 +++++++++++++ 5 files changed, 1230 insertions(+), 904 deletions(-) create mode 100644 catch/unit/memory/hipMemAdvise_old.cc diff --git a/catch/hipTestMain/config/config_amd_linux_common.json b/catch/hipTestMain/config/config_amd_linux_common.json index be119a27ab..a4e1fd0cb1 100644 --- a/catch/hipTestMain/config/config_amd_linux_common.json +++ b/catch/hipTestMain/config/config_amd_linux_common.json @@ -15,6 +15,8 @@ "Unit_hipMemset_Negative_OutOfBoundsPtr", "Unit_hipDeviceReset_Positive_Basic", "Unit_hipDeviceReset_Positive_Threaded", + "Unit_hipMemAdvise_AccessedBy_All_Devices", + "Unit_hipMemAdvise_No_Flag_Interference", "Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep", "Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep_ClonedGrph", "Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep_ChldNode", diff --git a/catch/hipTestMain/config/config_amd_windows_common.json b/catch/hipTestMain/config/config_amd_windows_common.json index 6bb11e842c..0e87e64b1f 100644 --- a/catch/hipTestMain/config/config_amd_windows_common.json +++ b/catch/hipTestMain/config/config_amd_windows_common.json @@ -94,6 +94,8 @@ "Unit_hipStreamSynchronize_NullStreamAndStreamPerThread", "Note: intermittent Seg fault failure ", "Unit_hipGraphAddEventRecordNode_Functional_WithoutFlags", + "Unit_hipMemAdvise_AccessedBy_All_Devices", + "Unit_hipMemAdvise_No_Flag_Interference", "Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep", "Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep_ClonedGrph", "Unit_hipGraphDestroyNode_Complx_ChkNumOfNodesNDep_ChldNode", diff --git a/catch/unit/memory/CMakeLists.txt b/catch/unit/memory/CMakeLists.txt index 876dae382b..2fa38dd802 100644 --- a/catch/unit/memory/CMakeLists.txt +++ b/catch/unit/memory/CMakeLists.txt @@ -108,6 +108,7 @@ set(TEST_SRC hipMemcpySync.cc hipMemsetSync.cc hipMemsetAsync.cc + hipMemAdvise_old.cc hipMemAdvise.cc hipMemRangeGetAttributes.cc hipStreamAttachMemAsync.cc @@ -194,6 +195,7 @@ set(TEST_SRC hipMemcpySync.cc hipMemsetSync.cc hipMemsetAsync.cc + hipMemAdvise_old.cc hipMemAdvise.cc hipMemRangeGetAttributes.cc hipMemRangeGetAttributes_old.cc diff --git a/catch/unit/memory/hipMemAdvise.cc b/catch/unit/memory/hipMemAdvise.cc index e5412ddfc7..87059a594e 100644 --- a/catch/unit/memory/hipMemAdvise.cc +++ b/catch/unit/memory/hipMemAdvise.cc @@ -1,943 +1,322 @@ /* -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 -IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER -LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -/* Test Case Description: - Scenario-1: The following Function Tests the working of flags which can be - assigned to HMM memory using hipMemAdvise() api - Scenario-2: Negative tests on hipMemAdvise() api - Scenario-3: The following function tests various scenarios around the flag - 'hipMemAdviseSetPreferredLocation' using HMM memory and hipMemAdvise() api - Scenario-4: The following function tests various scenarios around the flag - 'hipMemAdviseSetReadMostly' using HMM memory and hipMemAdvise() api - Scenario-5: The following function verifies if assigning of a flag - invalidates the earlier flag which was assigned to the same memory region - using hipMemAdvise() - Scenario-6: The following function tests if peers can set - hipMemAdviseSetAccessedBy flag - on HMM memory prefetched on each of the other gpus - Scenario-7: Set AccessedBy flag and check value returned by - hipMemRangeGetAttribute() It should be -2(same is observed on cuda) - Scenario-8: Set AccessedBy flag to device 0 on Hmm memory and prefetch the - memory to device 1, then probe for AccessedBy flag using - hipMemRangeGetAttribute() we should still see the said flag is set for - device 0 - Scenario-9: 1) Set AccessedBy to device 0 followed by PreferredLocation to - device 1 check for AccessedBy flag using hipMemRangeGetAttribute() it should - return 0 - 2) Unset AccessedBy to 0 and set it to device 1 followed by - PreferredLocation to device 1, check for AccessedBy flag using - hipMemRangeGetAttribute() it should return 1 - Scenario-10: Set AccessedBy flag to HMM memory launch a kernel and then unset - AccessedBy, launch kernel. We should not have any access issues - Scenario-11: Allocate memory using aligned_alloc(), assign PreferredLocation - flag to the allocated memory and launch a kernel. Kernel should get executed - successfully without hang or segfault - Scenario-12: Allocate Hmm memory, set advise to PreferredLocation and then - get attribute using the api hipMemRangeGetAttribute() for - hipMemRangeAttributeLastPrefetchLocation the value returned should be -2 - Scenario-13: Allocate HMM memory, set PreferredLocation to device 0, Prfetch - the mem to device1, probe for hipMemRangeAttributeLastPrefetchLocation using - hipMemRangeGetAttribute(), we should get 1 - Scenario-14: Allocate HMM memory, set ReadMostly followed by - PreferredLocation, probe for hipMemRangeAttributeReadMostly and - hipMemRangeAttributePreferredLocation - using hipMemRangeGetAttribute() we should observe 1 and 0 correspondingly. - In other words setting of hipMemRangeAttributePreferredLocation should not - impact hipMemRangeAttributeReadMostly advise to the memory - Scenario-15: Allocate Hmm memory, advise it to ReadMostly for gpu: 0 and - launch kernel on all other gpus except 0. This test case may discover any - effect or access denial case arising due to setting ReadMostly only to a - particular gpu -*/ - #include -#if __linux__ -#include -#include -#include -#endif +#include +#include +#include -// Kernel function -__global__ void MemAdvseKernel(int n, int *x) { - int index = blockIdx.x * blockDim.x + threadIdx.x; - if (index < n) - x[index] = x[index] * x[index]; -} - -// Kernel -__global__ void MemAdvise2(int *Hmm, int n) { - int index = blockIdx.x * blockDim.x + threadIdx.x; - int stride = blockDim.x * gridDim.x; - for (int i = index; i < n; i += stride) { - Hmm[i] = Hmm[i] + 10; +static inline hipMemoryAdvise GetUnsetMemAdvice(const hipMemoryAdvise advice) { + switch (advice) { + case hipMemAdviseSetAccessedBy: + return hipMemAdviseUnsetAccessedBy; + case hipMemAdviseSetReadMostly: + return hipMemAdviseUnsetReadMostly; + case hipMemAdviseSetPreferredLocation: + return hipMemAdviseUnsetPreferredLocation; + default: + assert("Invalid hipMemoryAdvise enumerator"); + return advice; } } -// Kernel -__global__ void MemAdvise3(int *Hmm, int *Hmm1, int n) { - int index = blockIdx.x * blockDim.x + threadIdx.x; - int stride = blockDim.x * gridDim.x; - for (int i = index; i < n; i += stride) { - Hmm1[i] = Hmm[i] + 10; +static inline hipMemRangeAttribute GetMemAdviceAttr(const hipMemoryAdvise advice) { + switch (advice) { + case hipMemAdviseSetAccessedBy: + return hipMemRangeAttributeAccessedBy; + case hipMemAdviseSetReadMostly: + return hipMemRangeAttributeReadMostly; + case hipMemAdviseSetPreferredLocation: + return hipMemRangeAttributePreferredLocation; + default: + assert("Invalid hipMemoryAdvise enumerator"); + return static_cast(-1); } } - -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 working of flags which can be assigned -// to HMM memory using hipMemAdvise() api -TEST_CASE("Unit_hipMemAdvise_TstFlags") { - int MangdMem = HmmAttrPrint(); - if (MangdMem == 1) { - bool IfTestPassed = true; - int NumDevs = 0, *Outpt = nullptr; - int MEM_SIZE = 4*1024, A_CONST = 9999; - float *Hmm = nullptr; - int AttrVal = 0; - HIP_CHECK(hipGetDeviceCount(&NumDevs)); - Outpt = new int[NumDevs]; - HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE * 2, hipMemAttachGlobal)); - // With the following for loop we iterate through each of the Gpus in the - // system set and unset the flags and check the behavior. - for (int i = 0; i < NumDevs; ++i) { - HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2, hipMemAdviseSetReadMostly, i)); - HIP_CHECK(hipMemRangeGetAttribute(&AttrVal, sizeof(AttrVal), - hipMemRangeAttributeReadMostly, Hmm, - MEM_SIZE * 2)); - if (AttrVal != 1) { - WARN("Attempt to set hipMemAdviseSetReadMostly flag failed!\n"); - IfTestPassed = false; - } - HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2, hipMemAdviseUnsetReadMostly, - i)); - - HIP_CHECK(hipMemRangeGetAttribute(&AttrVal, sizeof(AttrVal), - hipMemRangeAttributeReadMostly, Hmm, - (MEM_SIZE * 2))); - if (AttrVal != 0) { - WARN("Attempt to Unset hipMemAdviseSetReadMostly flag failed!\n"); - IfTestPassed = false; - } - AttrVal = A_CONST; - // Currently hipMemAdviseSetPreferredLocation and - // hipMemAdviseSetAccessedBy - // flags are resulting in issues: SWDEV-267357 - HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2, - hipMemAdviseSetPreferredLocation, i)); - HIP_CHECK(hipMemRangeGetAttribute(&AttrVal, sizeof(AttrVal), - hipMemRangeAttributePreferredLocation, - Hmm, (MEM_SIZE * 2))); - if (AttrVal != i) { - WARN("Attempt to set hipMemAdviseSetPreferredLocation flag failed!\n"); - IfTestPassed = false; - } - AttrVal = A_CONST; - HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2, - hipMemAdviseUnsetPreferredLocation, i)); - HIP_CHECK(hipMemRangeGetAttribute(&AttrVal, sizeof(AttrVal), - hipMemRangeAttributePreferredLocation, - Hmm, (MEM_SIZE * 2))); - if (AttrVal == i) { - WARN("Attempt to Unset hipMemAdviseUnsetPreferredLocation "); - WARN("flag failed!\n"); - IfTestPassed = false; - } - for (int m = 0; m < NumDevs; ++m) { - Outpt[m] = A_CONST; - } - HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2, hipMemAdviseSetAccessedBy, i)); - HIP_CHECK(hipMemRangeGetAttribute(Outpt, sizeof(Outpt), - hipMemRangeAttributeAccessedBy, Hmm, - (MEM_SIZE * 2))); - if ((Outpt[0]) != i) { - WARN("Attempt to set hipMemAdviseSetAccessedBy flag failed!\n"); - IfTestPassed = false; - } - for (int m = 0; m < NumDevs; ++m) { - Outpt[m] = A_CONST; - } - HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2, hipMemAdviseUnsetAccessedBy, - i)); - HIP_CHECK(hipMemRangeGetAttribute(Outpt, sizeof(Outpt), - hipMemRangeAttributeAccessedBy, Hmm, - (MEM_SIZE * 2))); - if ((Outpt[0]) >= 0) { - WARN("Attempt to Unset hipMemAdviseUnsetAccessedBy flag failed!\n"); - IfTestPassed = false; - } +std::vector GetDevicesWithAdviseSupport() { + const auto device_count = HipTest::getDeviceCount(); + std::vector supported_devices; + supported_devices.reserve(device_count + 1); + for (int i = 0; i < device_count; ++i) { + if (DeviceAttributesSupport(i, hipDeviceAttributeManagedMemory, + hipDeviceAttributeConcurrentManagedAccess)) { + supported_devices.push_back(i); + } + } + return supported_devices; +} + +TEST_CASE("Unit_hipMemAdvise_Set_Unset_Basic") { + auto supported_devices = GetDevicesWithAdviseSupport(); + if (supported_devices.empty()) { + HipTest::HIP_SKIP_TEST("Test needs at least 1 device that supports managed memory"); + return; + } + supported_devices.push_back(hipCpuDeviceId); + const auto device = GENERATE_COPY(from_range(supported_devices)); + + const auto SetUnset = [=](const hipMemoryAdvise advice) { + LinearAllocGuard alloc(LinearAllocs::hipMallocManaged, kPageSize); + int32_t attribute = 0; + HIP_CHECK(hipMemAdvise(alloc.ptr(), kPageSize, advice, device)); + HIP_CHECK(hipMemRangeGetAttribute(&attribute, sizeof(attribute), GetMemAdviceAttr(advice), + alloc.ptr(), kPageSize)); + REQUIRE((advice == hipMemAdviseSetReadMostly ? 1 : device) == attribute); + HIP_CHECK(hipMemAdvise(alloc.ptr(), kPageSize, GetUnsetMemAdvice(advice), device)); + HIP_CHECK(hipMemRangeGetAttribute(&attribute, sizeof(attribute), GetMemAdviceAttr(advice), + alloc.ptr(), kPageSize)); + REQUIRE((advice == hipMemAdviseSetReadMostly ? 0 : hipInvalidDeviceId) == attribute); + }; + +// Disabled due to defect - EXSWHTEC-132 +#if HT_NVIDIA + SECTION("hipMemAdviseSetAccessedBy") { SetUnset(hipMemAdviseSetAccessedBy); } +#endif + SECTION("hipMemAdviseSetReadMostly") { SetUnset(hipMemAdviseSetReadMostly); } + SECTION("hipMemAdviseSetPreferredLocation") { SetUnset(hipMemAdviseSetPreferredLocation); } +} + +TEST_CASE("Unit_hipMemAdvise_No_Flag_Interference") { + auto supported_devices = GetDevicesWithAdviseSupport(); + if (supported_devices.empty()) { + HipTest::HIP_SKIP_TEST("Test needs at least 1 device that supports managed memory"); + return; + } + supported_devices.push_back(hipCpuDeviceId); + const auto device = GENERATE_COPY(from_range(supported_devices)); + + std::array advice{hipMemAdviseSetReadMostly, hipMemAdviseSetPreferredLocation, + hipMemAdviseSetAccessedBy}; + for (int i = 0; i < 6; ++i) { + std::next_permutation(std::begin(advice), std::end(advice)); + LinearAllocGuard alloc(LinearAllocs::hipMallocManaged, kPageSize); + + for (const auto a : advice) { + HIP_CHECK(hipMemAdvise(alloc.ptr(), kPageSize, a, device)); + } + + for (const auto a : advice) { + int32_t attribute = 0; + HIP_CHECK(hipMemRangeGetAttribute(&attribute, sizeof(attribute), GetMemAdviceAttr(a), + alloc.ptr(), kPageSize)); + REQUIRE((a == hipMemAdviseSetReadMostly ? 1 : device) == attribute); } - delete [] Outpt; - HIP_CHECK(hipFree(Hmm)); - REQUIRE(IfTestPassed); - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); } } -TEST_CASE("Unit_hipMemAdvise_NegtveTsts") { - HipTest::HIP_SKIP_TEST("Fixed few issues to match with Nvidia, Skip now to avoid CI failures"); - return; - int MangdMem = HmmAttrPrint(); - if (MangdMem == 1) { - bool IfTestPassed = true; - int NumDevs = 0, MEM_SIZE = 4*1024; - float *Hmm = nullptr; - std::string str; - HIP_CHECK(hipGetDeviceCount(&NumDevs)); - HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE * 2, hipMemAttachGlobal)); +TEST_CASE("Unit_hipMemAdvise_Rounding") { + auto supported_devices = GetDevicesWithAdviseSupport(); + if (supported_devices.empty()) { + HipTest::HIP_SKIP_TEST("Test needs at least 1 device that supports managed memory"); + return; + } + supported_devices.push_back(hipCpuDeviceId); + const auto device = supported_devices.front(); + + LinearAllocGuard alloc(LinearAllocs::hipMallocManaged, 3 * kPageSize); + REQUIRE_FALSE(reinterpret_cast(alloc.ptr()) % kPageSize); + const auto [offset, width] = + GENERATE_COPY(std::make_pair(kPageSize / 4, kPageSize / 2), // Withing page + std::make_pair(kPageSize / 2, kPageSize), // Across page border + std::make_pair(kPageSize / 2, kPageSize * 2)); // Across two page borders + HIP_CHECK(hipMemAdvise(alloc.ptr() + offset, width, hipMemAdviseSetAccessedBy, device)); + constexpr auto RoundDown = [](const intptr_t a, const intptr_t n) { return a - a % n; }; + constexpr auto RoundUp = [RoundDown](const intptr_t a, const intptr_t n) { + return RoundDown(a + n - 1, n); + }; + const auto base = alloc.ptr(); + const auto rounded_up = RoundUp(offset + width, kPageSize); + unsigned int attribute = 0; + HIP_CHECK(hipMemRangeGetAttribute(&attribute, sizeof(attribute), hipMemRangeAttributeAccessedBy, + reinterpret_cast(base), rounded_up)); + REQUIRE(device == attribute); + HIP_CHECK(hipMemRangeGetAttribute(&attribute, sizeof(attribute), hipMemRangeAttributeAccessedBy, + alloc.ptr(), 3 * kPageSize)); + REQUIRE((rounded_up == 3 * kPageSize ? device : hipInvalidDeviceId) == attribute); +} + +TEST_CASE("Unit_hipMemAdvise_Flags_Do_Not_Cause_Prefetch") { + auto supported_devices = GetDevicesWithAdviseSupport(); + if (supported_devices.empty()) { + HipTest::HIP_SKIP_TEST("Test needs at least 1 device that supports managed memory"); + } + supported_devices.push_back(hipCpuDeviceId); + + const auto Test = [](const int device, const hipMemoryAdvise advice) { + LinearAllocGuard alloc(LinearAllocs::hipMallocManaged, kPageSize); + HIP_CHECK(hipMemAdvise(alloc.ptr(), kPageSize, advice, device)); + int32_t attribute = 0; + HIP_CHECK(hipMemRangeGetAttribute(&attribute, sizeof(attribute), + hipMemRangeAttributeLastPrefetchLocation, alloc.ptr(), + kPageSize)); + REQUIRE(attribute == hipInvalidDeviceId); + }; + const auto device = + GENERATE_COPY(from_range(std::begin(supported_devices), std::end(supported_devices))); + + SECTION("hipMemAdviseSetPreferredLocation") { Test(device, hipMemAdviseSetPreferredLocation); } +// Disabled on AMD due to defect - EXSWHTEC-132 +#if HT_NVIDIA + SECTION("hipMemAdviseSetAccessedBy") { Test(device, hipMemAdviseSetAccessedBy); } +#endif +} + +TEST_CASE("Unit_hipMemAdvise_Read_Write_After_Advise") { + auto supported_devices = GetDevicesWithAdviseSupport(); + if (supported_devices.empty()) { + HipTest::HIP_SKIP_TEST("Test needs at least 1 device that supports managed memory"); + } + LinearAllocGuard alloc(LinearAllocs::hipMallocManaged, kPageSize); + constexpr size_t count = kPageSize / sizeof(*alloc.ptr()); + + const auto ReadWriteManagedMemory = [&](const int device, const hipMemoryAdvise advice) { + HIP_CHECK(hipMemAdvise(alloc.ptr(), kPageSize, advice, device)); + + std::fill_n(alloc.ptr(), count, -1); + ArrayFindIfNot(alloc.ptr(), -1, count); + for (int i = 0; i < supported_devices.size(); ++i) { + HIP_CHECK(hipSetDevice(supported_devices[i])); + VectorIncrement<<>>(alloc.ptr(), 1, count); + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipDeviceSynchronize()); + ArrayFindIfNot(alloc.ptr(), i, count); + } + + int32_t attribute = 0; + HIP_CHECK(hipMemRangeGetAttribute(&attribute, sizeof(attribute), GetMemAdviceAttr(advice), + alloc.ptr(), kPageSize)); + REQUIRE((advice == hipMemAdviseSetReadMostly ? 1 : device) == attribute); + }; + +// Disabled on AMD due to defect - EXSWHTEC-133 +#if HT_NVIDIA + SECTION("ReadMostly") { ReadWriteManagedMemory(hipInvalidDeviceId, hipMemAdviseSetReadMostly); } +#endif + supported_devices.push_back(hipCpuDeviceId); + + const auto device = + GENERATE_COPY(from_range(std::begin(supported_devices), std::end(supported_devices))); + supported_devices.pop_back(); + SECTION("PreferredLocation") { ReadWriteManagedMemory(device, hipMemAdviseSetPreferredLocation); } + +// Disabled on AMD due to defect - EXSWHTEC-132 +#if HT_NVIDIA + SECTION("AccessedBy") { ReadWriteManagedMemory(device, hipMemAdviseSetAccessedBy); } +#endif +} + +TEST_CASE("Unit_hipMemAdvise_Prefetch_After_Advise") { + auto supported_devices = GetDevicesWithAdviseSupport(); + if (supported_devices.empty()) { + HipTest::HIP_SKIP_TEST("Test needs at least 1 device that supports managed memory"); + } + supported_devices.push_back(hipCpuDeviceId); + const auto advice = GENERATE(hipMemAdviseSetReadMostly, hipMemAdviseSetPreferredLocation + // Skipped due to defect - EXSWHTEC - 132 + // hipMemAdviseSetAccessedBy + ); + const auto device = GENERATE_COPY(from_range(supported_devices)); + + LinearAllocGuard alloc(LinearAllocs::hipMallocManaged, kPageSize); + HIP_CHECK(hipMemAdvise(alloc.ptr(), kPageSize, advice, device)); + + for (const auto d : supported_devices) { + HIP_CHECK(hipMemPrefetchAsync(alloc.ptr(), kPageSize, d)); + HIP_CHECK(hipStreamSynchronize(nullptr)); + int32_t attribute = 0; + HIP_CHECK(hipMemRangeGetAttribute(&attribute, sizeof(attribute), + hipMemRangeAttributeLastPrefetchLocation, alloc.ptr(), + kPageSize)); + REQUIRE(d == attribute); + } + + int32_t attribute = 0; + HIP_CHECK(hipMemRangeGetAttribute(&attribute, sizeof(attribute), GetMemAdviceAttr(advice), + alloc.ptr(), kPageSize)); + REQUIRE((advice == hipMemAdviseSetReadMostly ? 1 : device) == attribute); +} + +TEST_CASE("Unit_hipMemAdvise_AccessedBy_All_Devices") { + auto supported_devices = GetDevicesWithAdviseSupport(); + if (supported_devices.empty()) { + HipTest::HIP_SKIP_TEST("Test needs at least 1 device that supports managed memory"); + return; + } + supported_devices.push_back(hipCpuDeviceId); + + LinearAllocGuard alloc(LinearAllocs::hipMallocManaged, kPageSize); + for (const auto device : supported_devices) { + HIP_CHECK(hipMemAdvise(alloc.ptr(), kPageSize, hipMemAdviseSetAccessedBy, device)); + } + std::vector accessed_by(supported_devices.size(), hipInvalidDeviceId); + HIP_CHECK(hipMemRangeGetAttribute(accessed_by.data(), sizeof(accessed_by.data()), + hipMemRangeAttributeAccessedBy, alloc.ptr(), kPageSize)); + REQUIRE_THAT(accessed_by, Catch::Matchers::Equals(supported_devices)); +} + +TEST_CASE("Unit_hipMemAdvise_Negative_Parameters") { + auto supported_devices = GetDevicesWithAdviseSupport(); + if (supported_devices.empty()) { + HipTest::HIP_SKIP_TEST("Test needs at least 1 device that supports managed memory"); + } + const auto device = supported_devices.front(); + + LinearAllocGuard alloc(LinearAllocs::hipMallocManaged, kPageSize); + +// Disabled on NVIDIA due to defect - EXSWHTEC-122 #if HT_AMD - // Passing invalid value(99) device param - IfTestPassed &= CheckError(hipMemAdvise(Hmm, MEM_SIZE * 2, - hipMemAdviseSetReadMostly, 99), __LINE__); - - // Passing invalid value(-12) device param - IfTestPassed &= CheckError(hipMemAdvise(Hmm, MEM_SIZE * 2, - hipMemAdviseSetReadMostly, -12), __LINE__); -#endif - // Passing NULL as first parameter instead of valid pointer to a memory - IfTestPassed &= CheckError(hipMemAdvise(NULL, MEM_SIZE * 2, - hipMemAdviseSetReadMostly, 0), __LINE__); - - // Passing 0 for count(2nd param) parameter - IfTestPassed &= CheckError(hipMemAdvise(Hmm, 0, hipMemAdviseSetReadMostly, - 0), __LINE__); - - // Passing count much more than actually allocated value - IfTestPassed &= CheckError(hipMemAdvise(Hmm, MEM_SIZE * 6, - hipMemAdviseSetReadMostly, 0), __LINE__); - - REQUIRE(IfTestPassed); - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); + SECTION("Invalid advice") { + HIP_CHECK_ERROR(hipMemAdvise(alloc.ptr(), kPageSize, static_cast(-1), device), + hipErrorInvalidValue); } -} - -// The following function tests various scenarios around the flag -// 'hipMemAdviseSetPreferredLocation' using HMM memory and hipMemAdvise() api -TEST_CASE("Unit_hipMemAdvise_PrefrdLoc") { - int MangdMem = HmmAttrPrint(); - if (MangdMem == 1) { - // Check that when a page fault occurs for the memory region set to devPtr, - // the data is migrated to the destn processor - int MEM_SIZE = 4096, A_CONST = 9999; - int *Hmm = nullptr, NumDevs = 0, dev = A_CONST; - bool IfTestPassed = true; - HIP_CHECK(hipGetDeviceCount(&NumDevs)); - HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE * 3, hipMemAttachGlobal)); - for (int i = 0; i < ((MEM_SIZE * 3)/4); ++i) { - Hmm[i] = 4; - } - for (int devId = 0; devId < NumDevs; ++devId) { - HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE * 3, - hipMemAdviseSetPreferredLocation, devId)); - int NumElms = ((MEM_SIZE * 3)/4); - MemAdvseKernel<<>>(NumElms, Hmm); - int dev = A_CONST; - HIP_CHECK(hipMemRangeGetAttribute(&dev, sizeof(dev), - hipMemRangeAttributePreferredLocation, - Hmm, MEM_SIZE * 3)); - if (dev != devId) { - WARN("Memory observed to be not available on expected location\n"); - WARN("line no: " << __LINE__); - WARN("dev: " << dev); - IfTestPassed = false; - } - } - - // Check that when preferred location is set for a memory region, - // data can still be prefetched using hipMemPrefetchAsync - hipStream_t strm; - dev = A_CONST; - for (int devId = 0; devId < NumDevs; ++devId) { - HIP_CHECK(hipSetDevice(devId)); - HIP_CHECK(hipStreamCreate(&strm)); - HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE * 3, - hipMemAdviseSetPreferredLocation, devId)); - HIP_CHECK(hipMemPrefetchAsync(Hmm, MEM_SIZE * 3, devId, strm)); - HIP_CHECK(hipStreamSynchronize(strm)); - HIP_CHECK(hipMemRangeGetAttribute(&dev, sizeof(dev), - hipMemRangeAttributeLastPrefetchLocation, - Hmm, MEM_SIZE * 3)); - if (dev != devId) { - WARN("Memory reported to be not available at the Prefetched "); - WARN("location with device id: " << devId); - WARN("line no: " << __LINE__); - WARN("dev: " << dev); - IfTestPassed = false; - } - HIP_CHECK(hipStreamDestroy(strm)); - } - HIP_CHECK(hipFree(Hmm)); - REQUIRE(IfTestPassed); - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); - } -} - -// The following function tests various scenarios around the flag -// 'hipMemAdviseSetReadMostly' using HMM memory and hipMemAdvise() api - -TEST_CASE("Unit_hipMemAdvise_ReadMostly") { - int MangdMem = HmmAttrPrint(); - if (MangdMem == 1) { - bool IfTestPassed = true; - int MEM_SIZE = 4096, A_CONST = 9999; - float *Hmm = nullptr; - HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE)); - for (uint64_t i = 0; i < (MEM_SIZE/sizeof(float)); ++i) { - Hmm[i] = A_CONST; - } - HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetReadMostly, 0)); - // Checking if the data can be read after setting hipMemAdviseSetReadMostly - for (uint64_t i = 0; i < (MEM_SIZE/sizeof(float)); ++i) { - if (Hmm[i] != A_CONST) { - WARN("Didn't find expected value in Hmm memory after setting"); - WARN(" hipMemAdviseSetReadMostly flag line no.: " << __LINE__); - IfTestPassed = false; - } - } - - // Checking if the memory region can be modified - for (uint64_t i = 0; i < (MEM_SIZE/sizeof(float)); ++i) { - Hmm[i] = A_CONST; - } - - for (uint64_t i = 0; i < (MEM_SIZE/sizeof(float)); ++i) { - if (Hmm[i] != A_CONST) { - WARN("Didn't find expected value in Hmm memory after Modification\n"); - WARN("line no.: " << __LINE__); - IfTestPassed = false; - } - } - - int out = A_CONST; - HIP_CHECK(hipMemRangeGetAttribute(&out, 4, hipMemRangeAttributeReadMostly, - Hmm, MEM_SIZE)); - if (out != 1) { - WARN("out value: " << out); - IfTestPassed = false; - } - // Checking the advise attribute after prefetch - HIP_CHECK(hipMemPrefetchAsync(Hmm, MEM_SIZE, 0, 0)); - HIP_CHECK(hipDeviceSynchronize()); - HIP_CHECK(hipMemRangeGetAttribute(&out, sizeof(int), - hipMemRangeAttributeReadMostly, Hmm, - MEM_SIZE)); - if (out != 1) { - WARN("Attribute assigned to memory changed after calling "); - WARN("hipMemPrefetchAsync(). line no.: " << __LINE__); - WARN("out value: " << out); - IfTestPassed = false; - } - HIP_CHECK(hipFree(Hmm)); - REQUIRE(IfTestPassed); - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); - } -} - -// The following function verifies if assigning of a flag invalidates the -// earlier flag which was assigned to the same memory region using -// hipMemAdvise() -TEST_CASE("Unit_hipMemAdvise_TstFlgOverrideEffect") { - int MangdMem = HmmAttrPrint(); - if (MangdMem == 1) { - bool IfTestPassed = true; - int MEM_SIZE = 4*4096, A_CONST = 9999; - float *Hmm = nullptr; - int NumDevs = 0, dev = A_CONST; - - HIP_CHECK(hipGetDeviceCount(&NumDevs)); - HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE, hipMemAttachGlobal)); - for (int i = 0; i < NumDevs; ++i) { - HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetReadMostly, i)); - HIP_CHECK(hipMemRangeGetAttribute(&dev, sizeof(int), - hipMemRangeAttributeReadMostly, Hmm, - MEM_SIZE)); - if (dev != 1) { - WARN("hipMemAdviseSetReadMostly flag did not take affect despite "); - WARN("setting it using hipMemAdvise(). line no.: " << __LINE__); - IfTestPassed = false; - break; - } - dev = A_CONST; - HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetPreferredLocation, - i)); - HIP_CHECK(hipMemRangeGetAttribute(&dev, sizeof(int), - hipMemRangeAttributePreferredLocation, - Hmm, MEM_SIZE)); - if (dev != i) { - WARN("hipMemAdviseSetPreferredLocation flag did not take affect "); - WARN("despite setting it using hipMemAdvise()\n"); - WARN("line no.: " << __LINE__); - IfTestPassed = false; - break; - } - - HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetAccessedBy, i)); - dev = A_CONST; - HIP_CHECK(hipMemRangeGetAttribute(&dev, sizeof(int), - hipMemRangeAttributeAccessedBy, Hmm, - MEM_SIZE)); - if (dev != i) { - WARN("hipMemAdviseSetAccessedBy flag did not take affect despite "); - WARN("setting it using hipMemAdvise(). line no.: " << __LINE__); - IfTestPassed = false; - break; - } - HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseUnsetAccessedBy, i)); - } - HIP_CHECK(hipFree(Hmm)); - REQUIRE(IfTestPassed); - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); - } -} - - -// The following function tests if peers can set hipMemAdviseSetAccessedBy flag -// on HMM memory prefetched on each of the other gpus -#if HT_AMD -TEST_CASE("Unit_hipMemAdvise_TstAccessedByPeer") { - int MangdMem = HmmAttrPrint(); - if (MangdMem == 1) { - bool IfTestPassed = true; - int *Hmm = nullptr, MEM_SIZE = 4*4096, A_CONST = 9999;; - int NumDevs = 0, CanAccessPeer = A_CONST, flag = 0; - - HIP_CHECK(hipGetDeviceCount(&NumDevs)); - if (NumDevs < 2) { - SUCCEED("Test TestSetAccessedByPeer() need atleast two Gpus to test" - " the scenario. This system has GPUs less than 2"); - } - HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE, hipMemAttachGlobal)); - for (int i = 0; i < NumDevs; ++i) { - HIP_CHECK(hipMemPrefetchAsync(Hmm, MEM_SIZE, i, 0)); - for (int j = 0; j < NumDevs; ++j) { - if (i == j) - continue; - HIP_CHECK(hipSetDevice(j)); - HIP_CHECK(hipDeviceCanAccessPeer(&CanAccessPeer, j, i)); - if (CanAccessPeer) { - HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetAccessedBy, j)); - for (uint64_t m = 0; m < (MEM_SIZE/sizeof(int)); ++m) { - Hmm[m] = 4; - } - HIP_CHECK(hipDeviceEnablePeerAccess(i, 0)); - MemAdvseKernel<<<(MEM_SIZE/sizeof(int)/32), 32>>>( - (MEM_SIZE/sizeof(int)), Hmm); - HIP_CHECK(hipDeviceSynchronize()); - // Verifying the result - for (uint64_t m = 0; m < (MEM_SIZE/sizeof(int)); ++m) { - if (Hmm[m] != 16) { - flag = 1; - } - } - if (flag) { - WARN("Didnt get Expected results with device: " << j); - WARN("line no.: " << __LINE__); - IfTestPassed = false; - flag = 0; - } - } - } - } - HIP_CHECK(hipFree(Hmm)); - REQUIRE(IfTestPassed); - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); - } -} #endif - -/* Set AccessedBy flag and check value returned by hipMemRangeGetAttribute() - It should be -2(same is observed on cuda)*/ -TEST_CASE("Unit_hipMemAdvise_TstAccessedByFlg") { - int managed = HmmAttrPrint(); - if (managed == 1) { - int *Hmm = NULL, data = 999; - HIP_CHECK(hipMallocManaged(&Hmm, 2*4096)); - HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetAccessedBy, 0)); - HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), - hipMemRangeAttributeLastPrefetchLocation, - Hmm, 2*4096)); - if (data != -2) { - WARN("Didnt get expected value!!\n"); - REQUIRE(false); - } - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); - } -} - -/* Set AccessedBy flag to device 0 on Hmm memory and prefetch the memory to - device 1, then probe for AccessedBy flag using hipMemRangeGetAttribute() - we should still see the said flag is set for device 0*/ -TEST_CASE("Unit_hipMemAdvise_TstAccessedByFlg2") { - int managed = HmmAttrPrint(); - if (managed == 1) { - int *Hmm = NULL, data = 999, Ngpus = 0; - HIP_CHECK(hipGetDeviceCount(&Ngpus)); - if (Ngpus >= 2) { - hipStream_t strm; - HIP_CHECK(hipStreamCreate(&strm)); - HIP_CHECK(hipMallocManaged(&Hmm, 2*4096)); - HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetAccessedBy, 0)); - HIP_CHECK(hipMemPrefetchAsync(Hmm, 2*4096, 1, strm)); - HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), - hipMemRangeAttributeAccessedBy, Hmm, 2*4096)); - if (data != 0) { - WARN("Didnt get expected behavior at line: " << __LINE__); - REQUIRE(false); - } - HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseUnsetAccessedBy, 0)); - HIP_CHECK(hipStreamDestroy(strm)); - HIP_CHECK(hipFree(Hmm)); - } - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); - } -} - - - -/* 1) Set AccessedBy to device 0 followed by PreferredLocation to device 1 - check for AccessedBy flag using hipMemRangeGetAttribute() it should - return 0 - 2) Unset AccessedBy to 0 and set it to device 1 followed by - PreferredLocation to device 1, check for AccessedBy flag using - hipMemRangeGetAttribute() it should return 1*/ - -TEST_CASE("Unit_hipMemAdvise_TstAccessedByFlg3") { - int managed = HmmAttrPrint(); - if (managed == 1) { - int *Hmm = NULL, data = 999, Ngpus = 0; - HIP_CHECK(hipGetDeviceCount(&Ngpus)); - if (Ngpus >= 2) { - HIP_CHECK(hipMallocManaged(&Hmm, 2*4096)); - HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetAccessedBy, 0)); - HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetPreferredLocation, 1)); - HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), - hipMemRangeAttributeAccessedBy, Hmm, 2*4096)); - if (data != 0) { - WARN("Didnt get expected behavior at line: " << __LINE__); - REQUIRE(false); - } - HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseUnsetAccessedBy, 0)); - HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetAccessedBy, 1)); - HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetPreferredLocation, 0)); - HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), - hipMemRangeAttributeAccessedBy, Hmm, 2*4096)); - if (data != 1) { - WARN("Didnt get expected behavior at line: " << __LINE__); - REQUIRE(false); - } - HIP_CHECK(hipFree(Hmm)); - } - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); - } -} - - -/* Set AccessedBy flag to HMM memory launch a kernel and then unset - AccessedBy, launch kernel. We should not have any access issues*/ - -TEST_CASE("Unit_hipMemAdvise_TstAccessedByFlg4") { - int managed = HmmAttrPrint(); - if (managed == 1) { - int *Hmm = NULL, NumElms = (1024 * 1024), InitVal = 123, blockSize = 64; - int DataMismatch = 0; - hipStream_t strm; - HIP_CHECK(hipStreamCreate(&strm)); - HIP_CHECK(hipMallocManaged(&Hmm, (NumElms * sizeof(int)))); - HIP_CHECK(hipMemAdvise(Hmm, (NumElms * sizeof(int)), - hipMemAdviseSetAccessedBy, 0)); - // Initializing memory - for (int i = 0; i < NumElms; ++i) { - Hmm[i] = InitVal; - } - dim3 dimBlock(blockSize, 1, 1); - dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1); - // launching kernel from each one of the gpus - MemAdvise2<<>>(Hmm, NumElms); - HIP_CHECK(hipStreamSynchronize(strm)); - - // verifying the final result - for (int i = 0; i < NumElms; ++i) { - if (Hmm[i] != (InitVal + 10)) { - DataMismatch++; - } - } - - if (DataMismatch != 0) { - WARN("DataMismatch is observed at line: " << __LINE__); - REQUIRE(false); - } - - HIP_CHECK(hipMemAdvise(Hmm, (NumElms * sizeof(int)), - hipMemAdviseUnsetAccessedBy, 0)); - MemAdvise2<<>>(Hmm, NumElms); - HIP_CHECK(hipStreamSynchronize(strm)); - // verifying the final result - for (int i = 0; i < NumElms; ++i) { - if (Hmm[i] != (InitVal + (2*10))) { - DataMismatch++; - } - } - - if (DataMismatch != 0) { - WARN("DataMismatch is observed at line: " << __LINE__); - REQUIRE(false); - } - HIP_CHECK(hipFree(Hmm)); - HIP_CHECK(hipStreamDestroy(strm)); - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); - } -} - - -/* Allocate memory using aligned_alloc(), assign PreferredLocation flag to - the allocated memory and launch a kernel. Kernel should get executed - successfully without hang or segfault*/ -#if __linux__ && HT_AMD -TEST_CASE("Unit_hipMemAdvise_TstAlignedAllocMem") { - if ((setenv("HSA_XNACK", "1", 1)) != 0) { - WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!"); - REQUIRE(false); - } - // The following code block checks for gfx90a so as to skip if the device is not MI200 - - hipDeviceProp_t prop; - int device; - HIP_CHECK(hipGetDevice(&device)); - HIP_CHECK(hipGetDeviceProperties(&prop, device)); - std::string gfxName(prop.gcnArchName); - - if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) { - int stat = 0; - if (fork() == 0) { - // The below part should be inside fork - int managedMem = 0, pageMemAccess = 0; - HIP_CHECK(hipDeviceGetAttribute(&pageMemAccess, - hipDeviceAttributePageableMemoryAccess, 0)); - WARN("hipDeviceAttributePageableMemoryAccess:" << pageMemAccess); - - HIP_CHECK(hipDeviceGetAttribute(&managedMem, hipDeviceAttributeManagedMemory, 0)); - WARN("hipDeviceAttributeManagedMemory: " << managedMem); - if ((managedMem == 1) && (pageMemAccess == 1)) { - int *Mllc = nullptr, MemSz = 4096 * 4, NumElms = 4096, InitVal = 123; - // Mllc = reinterpret_cast<(int *)>(aligned_alloc(4096, MemSz)); - Mllc = reinterpret_cast(aligned_alloc(4096, 4096*4)); - for (int i = 0; i < NumElms; ++i) { - Mllc[i] = InitVal; - } - hipStream_t strm; - int DataMismatch = 0; - HIP_CHECK(hipStreamCreate(&strm)); - // The following hipMemAdvise() call is made to know if advise on - // aligned_alloc() is causing any issue - HIP_CHECK(hipMemAdvise(Mllc, MemSz, hipMemAdviseSetPreferredLocation, 0)); - HIP_CHECK(hipMemPrefetchAsync(Mllc, MemSz, 0, strm)); - HIP_CHECK(hipStreamSynchronize(strm)); - MemAdvise2<<<(NumElms/32), 32, 0, strm>>>(Mllc, NumElms); - HIP_CHECK(hipStreamSynchronize(strm)); - for (int i = 0; i < NumElms; ++i) { - if (Mllc[i] != (InitVal + 10)) { - DataMismatch++; - } - } - if (DataMismatch != 0) { - WARN("DataMismatch observed!!"); - exit(9); // 9 for failure - } else { - exit(10); // 10 for Pass result - } - } else { - SUCCEED("GPU 0 doesn't support ManagedMemory with hipDeviceAttributePageableMemoryAccess " - "attribute. Hence skipping the testing with Pass result.\n"); - exit(Catch::ResultDisposition::ContinueOnFailure); - } - } else { - wait(&stat); - int Result = WEXITSTATUS(stat); - if (Result == Catch::ResultDisposition::ContinueOnFailure) { - WARN("GPU 0 doesn't support ManagedMemory with hipDeviceAttributePageableMemoryAccess " - "attribute. Hence skipping the testing with Pass result.\n"); - } else { - if (Result != 10) { - REQUIRE(false); - } - } - } - } else { - SUCCEED("Memory model feature is only supported for gfx90a, Hence" - "skipping the testcase for this GPU " << device); - WARN("Memory model feature is only supported for gfx90a, Hence" - "skipping the testcase for this GPU " << device); + const auto advice = GENERATE(hipMemAdviseSetAccessedBy, hipMemAdviseSetReadMostly, + hipMemAdviseSetPreferredLocation); + SECTION("count == 0") { + HIP_CHECK_ERROR(hipMemAdvise(alloc.ptr(), 0, advice, device), hipErrorInvalidValue); } -} +// Disabled due to defect - EXSWHTEC-131 +#if HT_NVIDIA + SECTION("count larger than allocation size") { + HIP_CHECK_ERROR(hipMemAdvise(alloc.ptr(), kPageSize + 1, advice, device), hipErrorInvalidValue); + } #endif -/* Allocate Hmm memory, set advise to PreferredLocation and then get - attribute using the api hipMemRangeGetAttribute() for - hipMemRangeAttributeLastPrefetchLocation the value returned should be -2*/ - -TEST_CASE("Unit_hipMemAdvise_TstMemAdvisePrefrdLoc") { - int managed = HmmAttrPrint(); - if (managed == 1) { - int *Hmm = NULL, data = 999; - HIP_CHECK(hipMallocManaged(&Hmm, 4096)); - HIP_CHECK(hipMemAdvise(Hmm, 4096, hipMemAdviseSetPreferredLocation, 0)); - HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), - hipMemRangeAttributeLastPrefetchLocation, - Hmm, 4096)); - if (data != -2) { - WARN("Didnt receive expected value."); - REQUIRE(false); - } - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); + SECTION("dev_ptr == nullptr") { + HIP_CHECK_ERROR(hipMemAdvise(nullptr, kPageSize, advice, device), hipErrorInvalidValue); } -} - -/* Allocate HMM memory, set PreferredLocation to device 0, Prfetch the mem - to device1, probe for hipMemRangeAttributeLastPrefetchLocation using - hipMemRangeGetAttribute(), we should get 1*/ - -TEST_CASE("Unit_hipMemAdvise_TstMemAdviseLstPreftchLoc") { - int NumDevs = 0; - HIP_CHECK(hipGetDeviceCount(&NumDevs)); - if (NumDevs >= 2) { - int managed = HmmAttrPrint(); - if (managed == 1) { - int *Hmm = NULL, data = 999; - hipStream_t strm; - HIP_CHECK(hipSetDevice(1)); - HIP_CHECK(hipStreamCreate(&strm)); - HIP_CHECK(hipMallocManaged(&Hmm, 4096)); - HIP_CHECK(hipMemAdvise(Hmm, 4096, hipMemAdviseSetPreferredLocation, 0)); - HIP_CHECK(hipMemPrefetchAsync(Hmm, 4096, 1, strm)); - HIP_CHECK(hipStreamSynchronize(strm)); - HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), - hipMemRangeAttributeLastPrefetchLocation, - Hmm, 4096)); - if (data != 1) { - WARN("Didnt receive expected value!!"); - REQUIRE(false); - } - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); - } - } else { - SUCCEED("This system has less than 2 gpus hence skipping the test.\n"); + SECTION("dev_ptr pointing to non-managed memory") { + LinearAllocGuard alloc(LinearAllocs::hipMalloc, kPageSize); + HIP_CHECK_ERROR(hipMemAdvise(alloc.ptr(), kPageSize, advice, device), hipErrorInvalidValue); } -} - -/* Allocate HMM memory, set ReadMostly followed by PreferredLocation, probe - for hipMemRangeAttributeReadMostly and hipMemRangeAttributePreferredLocation - using hipMemRangeGetAttribute() we should observe 1 and 0 correspondingly. - In other words setting of hipMemRangeAttributePreferredLocation should not - impact hipMemRangeAttributeReadMostly advise to the memory*/ - -TEST_CASE("Unit_hipMemAdvise_TstMemAdviseMultiFlag") { - int managed = HmmAttrPrint(); - if (managed == 1) { - int *Hmm = NULL, data = 999; - HIP_CHECK(hipMallocManaged(&Hmm, 4096)); - HIP_CHECK(hipMemAdvise(Hmm, 4096, hipMemAdviseSetReadMostly, 0)); - HIP_CHECK(hipMemAdvise(Hmm, 4096, hipMemAdviseSetPreferredLocation, 0)); - HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), - hipMemRangeAttributeReadMostly, Hmm, - 4096)); - if (data != 1) { - WARN("Didnt receive expected value at line: " << data); - REQUIRE(false); - } - HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), - hipMemRangeAttributePreferredLocation, Hmm, - 4096)); - if (data != 0) { - WARN("Didnt receive expected value at line: " << data); - REQUIRE(false); - } - HIP_CHECK(hipFree(Hmm)); - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); +// Disabled on AMD due to defect - EXSWHTEC-130 +#if HT_NVIDIA + SECTION("Invalid device") { + HIP_CHECK_ERROR(hipMemAdvise(alloc.ptr(), kPageSize, advice, hipInvalidDeviceId), + (advice == hipMemAdviseSetReadMostly ? hipSuccess : hipErrorInvalidDevice)); } -} - - - -/*Allocate Hmm memory, advise it to ReadMostly for gpu: 0 and launch kernel - on all other gpus except 0. This test case may discover any effect or - access denial case arising due to setting ReadMostly only to a particular - gpu*/ - -TEST_CASE("Unit_hipMemAdvise_ReadMosltyMgpuTst") { - int managed = HmmAttrPrint(); - if (managed == 1) { - int Ngpus = 0; - HIP_CHECK(hipGetDeviceCount(&Ngpus)); - if (Ngpus < 2) { - SUCCEED("This test needs atleast two gpus to run." - "Hence skipping the test.\n"); - } - int *Hmm = NULL, NumElms = (1024 * 1024), InitVal = 123, blockSize = 64; - int *Hmm1 = NULL, DataMismatch = 0; - hipStream_t strm; - HIP_CHECK(hipStreamCreate(&strm)); - HIP_CHECK(hipMallocManaged(&Hmm, (NumElms * sizeof(int)))); - // Initializing memory - for (int i = 0; i < NumElms; ++i) { - Hmm[i] = InitVal; - } - HIP_CHECK(hipMemAdvise(Hmm, (NumElms * sizeof(int)), - hipMemAdviseSetReadMostly, 0)); - dim3 dimBlock(blockSize, 1, 1); - dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1); -#if HT_AMD - SECTION("Launch Kernel on all other gpus") { - // launching kernel from each one of the gpus - for (int i = 1; i < Ngpus; ++i) { - DataMismatch = 0; - HIP_CHECK(hipSetDevice(i)); - HIP_CHECK(hipMallocManaged(&Hmm1, (NumElms * sizeof(int)))); - MemAdvise3<<>>(Hmm, Hmm1, NumElms); - HIP_CHECK(hipStreamSynchronize(strm)); - // verifying the results - for (int j = 0; j < NumElms; ++j) { - if (Hmm1[j] != (InitVal + 10)) { - DataMismatch++; - } - } - if (DataMismatch != 0) { - WARN("DataMismatch is observed with the gpu: " << i); - REQUIRE(false); - } - HIP_CHECK(hipFree(Hmm1)); - } - } - - SECTION("Launch Kernel on all other gpus and manipulate the content") { - for (int i = 0; i < Ngpus; ++i) { - DataMismatch = 0; - HIP_CHECK(hipSetDevice(i)); - HIP_CHECK(hipMemAdvise(Hmm, (NumElms * sizeof(int)), - hipMemAdviseSetReadMostly, i)); - MemAdvise2<<>>(Hmm, NumElms); - HIP_CHECK(hipStreamSynchronize(strm)); - } - // verifying the final result - for (int i = 0; i < NumElms; ++i) { - if (Hmm[i] != (InitVal + Ngpus * 10)) { - DataMismatch++; - } - } - - if (DataMismatch != 0) { - WARN("DataMismatch is observed at line: " << __LINE__); - REQUIRE(false); - } - } #endif - HIP_CHECK(hipFree(Hmm)); - HIP_CHECK(hipStreamDestroy(strm)); - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); - } } - - -TEST_CASE("Unit_hipMemAdvise_TstSetUnsetPrfrdLoc") { - int managed = HmmAttrPrint(); - if (managed == 1) { - int *Hmm = NULL, data = 999; - HIP_CHECK(hipMallocManaged(&Hmm, 2*4096)); - HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetPreferredLocation, 0)); - HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), - hipMemRangeAttributePreferredLocation, Hmm, 2*4096)); - if (data != 0) { - WARN("Didnt receive expected value!!"); - REQUIRE(false); - } - HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseUnsetPreferredLocation, 0)); - HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), - hipMemRangeAttributePreferredLocation, Hmm, 2*4096)); - if (data != -2) { - WARN("Didnt receive expected value!!"); - REQUIRE(false); - } - } else { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " - "attribute. Hence skipping the testing with Pass result.\n"); - } -} - diff --git a/catch/unit/memory/hipMemAdvise_old.cc b/catch/unit/memory/hipMemAdvise_old.cc new file mode 100644 index 0000000000..cef164e2ff --- /dev/null +++ b/catch/unit/memory/hipMemAdvise_old.cc @@ -0,0 +1,941 @@ +/* +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, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/* Test Case Description: + Scenario-1: The following Function Tests the working of flags which can be + assigned to HMM memory using hipMemAdvise() api + Scenario-2: Negative tests on hipMemAdvise() api + Scenario-3: The following function tests various scenarios around the flag + 'hipMemAdviseSetPreferredLocation' using HMM memory and hipMemAdvise() api + Scenario-4: The following function tests various scenarios around the flag + 'hipMemAdviseSetReadMostly' using HMM memory and hipMemAdvise() api + Scenario-5: The following function verifies if assigning of a flag + invalidates the earlier flag which was assigned to the same memory region + using hipMemAdvise() + Scenario-6: The following function tests if peers can set + hipMemAdviseSetAccessedBy flag + on HMM memory prefetched on each of the other gpus + Scenario-7: Set AccessedBy flag and check value returned by + hipMemRangeGetAttribute() It should be -2(same is observed on cuda) + Scenario-8: Set AccessedBy flag to device 0 on Hmm memory and prefetch the + memory to device 1, then probe for AccessedBy flag using + hipMemRangeGetAttribute() we should still see the said flag is set for + device 0 + Scenario-9: 1) Set AccessedBy to device 0 followed by PreferredLocation to + device 1 check for AccessedBy flag using hipMemRangeGetAttribute() it should + return 0 + 2) Unset AccessedBy to 0 and set it to device 1 followed by + PreferredLocation to device 1, check for AccessedBy flag using + hipMemRangeGetAttribute() it should return 1 + Scenario-10: Set AccessedBy flag to HMM memory launch a kernel and then unset + AccessedBy, launch kernel. We should not have any access issues + Scenario-11: Allocate memory using aligned_alloc(), assign PreferredLocation + flag to the allocated memory and launch a kernel. Kernel should get executed + successfully without hang or segfault + Scenario-12: Allocate Hmm memory, set advise to PreferredLocation and then + get attribute using the api hipMemRangeGetAttribute() for + hipMemRangeAttributeLastPrefetchLocation the value returned should be -2 + Scenario-13: Allocate HMM memory, set PreferredLocation to device 0, Prfetch + the mem to device1, probe for hipMemRangeAttributeLastPrefetchLocation using + hipMemRangeGetAttribute(), we should get 1 + Scenario-14: Allocate HMM memory, set ReadMostly followed by + PreferredLocation, probe for hipMemRangeAttributeReadMostly and + hipMemRangeAttributePreferredLocation + using hipMemRangeGetAttribute() we should observe 1 and 0 correspondingly. + In other words setting of hipMemRangeAttributePreferredLocation should not + impact hipMemRangeAttributeReadMostly advise to the memory + Scenario-15: Allocate Hmm memory, advise it to ReadMostly for gpu: 0 and + launch kernel on all other gpus except 0. This test case may discover any + effect or access denial case arising due to setting ReadMostly only to a + particular gpu +*/ + +#include +#if __linux__ +#include +#include +#include +#endif + +// Kernel function +__global__ void MemAdvseKernel(int n, int *x) { + int index = blockIdx.x * blockDim.x + threadIdx.x; + if (index < n) + x[index] = x[index] * x[index]; +} + +// Kernel +__global__ void MemAdvise2(int *Hmm, int n) { + int index = blockIdx.x * blockDim.x + threadIdx.x; + int stride = blockDim.x * gridDim.x; + for (int i = index; i < n; i += stride) { + Hmm[i] = Hmm[i] + 10; + } +} + +// Kernel +__global__ void MemAdvise3(int *Hmm, int *Hmm1, int n) { + int index = blockIdx.x * blockDim.x + threadIdx.x; + int stride = blockDim.x * gridDim.x; + for (int i = index; i < n; i += stride) { + Hmm1[i] = Hmm[i] + 10; + } +} + + +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 working of flags which can be assigned +// to HMM memory using hipMemAdvise() api +TEST_CASE("Unit_hipMemAdvise_TstFlags") { + int MangdMem = HmmAttrPrint(); + if (MangdMem == 1) { + bool IfTestPassed = true; + int NumDevs = 0, *Outpt = nullptr; + int MEM_SIZE = 4*1024, A_CONST = 9999; + float *Hmm = nullptr; + int AttrVal = 0; + HIP_CHECK(hipGetDeviceCount(&NumDevs)); + Outpt = new int[NumDevs]; + HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE * 2, hipMemAttachGlobal)); + // With the following for loop we iterate through each of the Gpus in the + // system set and unset the flags and check the behavior. + for (int i = 0; i < NumDevs; ++i) { + HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2, hipMemAdviseSetReadMostly, i)); + HIP_CHECK(hipMemRangeGetAttribute(&AttrVal, sizeof(AttrVal), + hipMemRangeAttributeReadMostly, Hmm, + MEM_SIZE * 2)); + if (AttrVal != 1) { + WARN("Attempt to set hipMemAdviseSetReadMostly flag failed!\n"); + IfTestPassed = false; + } + HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2, hipMemAdviseUnsetReadMostly, + i)); + + HIP_CHECK(hipMemRangeGetAttribute(&AttrVal, sizeof(AttrVal), + hipMemRangeAttributeReadMostly, Hmm, + (MEM_SIZE * 2))); + if (AttrVal != 0) { + WARN("Attempt to Unset hipMemAdviseSetReadMostly flag failed!\n"); + IfTestPassed = false; + } + AttrVal = A_CONST; + // Currently hipMemAdviseSetPreferredLocation and + // hipMemAdviseSetAccessedBy + // flags are resulting in issues: SWDEV-267357 + HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2, + hipMemAdviseSetPreferredLocation, i)); + HIP_CHECK(hipMemRangeGetAttribute(&AttrVal, sizeof(AttrVal), + hipMemRangeAttributePreferredLocation, + Hmm, (MEM_SIZE * 2))); + if (AttrVal != i) { + WARN("Attempt to set hipMemAdviseSetPreferredLocation flag failed!\n"); + IfTestPassed = false; + } + AttrVal = A_CONST; + HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2, + hipMemAdviseUnsetPreferredLocation, i)); + HIP_CHECK(hipMemRangeGetAttribute(&AttrVal, sizeof(AttrVal), + hipMemRangeAttributePreferredLocation, + Hmm, (MEM_SIZE * 2))); + if (AttrVal == i) { + WARN("Attempt to Unset hipMemAdviseUnsetPreferredLocation "); + WARN("flag failed!\n"); + IfTestPassed = false; + } + for (int m = 0; m < NumDevs; ++m) { + Outpt[m] = A_CONST; + } + HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2, hipMemAdviseSetAccessedBy, i)); + HIP_CHECK(hipMemRangeGetAttribute(Outpt, sizeof(Outpt), + hipMemRangeAttributeAccessedBy, Hmm, + (MEM_SIZE * 2))); + if ((Outpt[0]) != i) { + WARN("Attempt to set hipMemAdviseSetAccessedBy flag failed!\n"); + IfTestPassed = false; + } + for (int m = 0; m < NumDevs; ++m) { + Outpt[m] = A_CONST; + } + HIP_CHECK(hipMemAdvise(Hmm , MEM_SIZE * 2, hipMemAdviseUnsetAccessedBy, + i)); + HIP_CHECK(hipMemRangeGetAttribute(Outpt, sizeof(Outpt), + hipMemRangeAttributeAccessedBy, Hmm, + (MEM_SIZE * 2))); + if ((Outpt[0]) >= 0) { + WARN("Attempt to Unset hipMemAdviseUnsetAccessedBy flag failed!\n"); + IfTestPassed = false; + } + } + delete [] Outpt; + HIP_CHECK(hipFree(Hmm)); + REQUIRE(IfTestPassed); + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + +TEST_CASE("Unit_hipMemAdvise_NegtveTsts") { + int MangdMem = HmmAttrPrint(); + if (MangdMem == 1) { + bool IfTestPassed = true; + int NumDevs = 0, MEM_SIZE = 4*1024; + float *Hmm = nullptr; + std::string str; + HIP_CHECK(hipGetDeviceCount(&NumDevs)); + HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE * 2, hipMemAttachGlobal)); +#if HT_AMD + // Passing invalid value(99) device param + IfTestPassed &= CheckError(hipMemAdvise(Hmm, MEM_SIZE * 2, + hipMemAdviseSetReadMostly, 99), __LINE__); + + // Passing invalid value(-12) device param + IfTestPassed &= CheckError(hipMemAdvise(Hmm, MEM_SIZE * 2, + hipMemAdviseSetReadMostly, -12), __LINE__); +#endif + // Passing NULL as first parameter instead of valid pointer to a memory + IfTestPassed &= CheckError(hipMemAdvise(NULL, MEM_SIZE * 2, + hipMemAdviseSetReadMostly, 0), __LINE__); + + // Passing 0 for count(2nd param) parameter + IfTestPassed &= CheckError(hipMemAdvise(Hmm, 0, hipMemAdviseSetReadMostly, + 0), __LINE__); + + // Passing count much more than actually allocated value + IfTestPassed &= CheckError(hipMemAdvise(Hmm, MEM_SIZE * 6, + hipMemAdviseSetReadMostly, 0), __LINE__); + + REQUIRE(IfTestPassed); + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + +// The following function tests various scenarios around the flag +// 'hipMemAdviseSetPreferredLocation' using HMM memory and hipMemAdvise() api +TEST_CASE("Unit_hipMemAdvise_PrefrdLoc") { + int MangdMem = HmmAttrPrint(); + if (MangdMem == 1) { + // Check that when a page fault occurs for the memory region set to devPtr, + // the data is migrated to the destn processor + int MEM_SIZE = 4096, A_CONST = 9999; + int *Hmm = nullptr, NumDevs = 0, dev = A_CONST; + bool IfTestPassed = true; + HIP_CHECK(hipGetDeviceCount(&NumDevs)); + HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE * 3, hipMemAttachGlobal)); + for (int i = 0; i < ((MEM_SIZE * 3)/4); ++i) { + Hmm[i] = 4; + } + for (int devId = 0; devId < NumDevs; ++devId) { + HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE * 3, + hipMemAdviseSetPreferredLocation, devId)); + int NumElms = ((MEM_SIZE * 3)/4); + MemAdvseKernel<<>>(NumElms, Hmm); + int dev = A_CONST; + HIP_CHECK(hipMemRangeGetAttribute(&dev, sizeof(dev), + hipMemRangeAttributePreferredLocation, + Hmm, MEM_SIZE * 3)); + if (dev != devId) { + WARN("Memory observed to be not available on expected location\n"); + WARN("line no: " << __LINE__); + WARN("dev: " << dev); + IfTestPassed = false; + } + } + + // Check that when preferred location is set for a memory region, + // data can still be prefetched using hipMemPrefetchAsync + hipStream_t strm; + dev = A_CONST; + for (int devId = 0; devId < NumDevs; ++devId) { + HIP_CHECK(hipSetDevice(devId)); + HIP_CHECK(hipStreamCreate(&strm)); + HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE * 3, + hipMemAdviseSetPreferredLocation, devId)); + HIP_CHECK(hipMemPrefetchAsync(Hmm, MEM_SIZE * 3, devId, strm)); + HIP_CHECK(hipStreamSynchronize(strm)); + HIP_CHECK(hipMemRangeGetAttribute(&dev, sizeof(dev), + hipMemRangeAttributeLastPrefetchLocation, + Hmm, MEM_SIZE * 3)); + if (dev != devId) { + WARN("Memory reported to be not available at the Prefetched "); + WARN("location with device id: " << devId); + WARN("line no: " << __LINE__); + WARN("dev: " << dev); + IfTestPassed = false; + } + HIP_CHECK(hipStreamDestroy(strm)); + } + HIP_CHECK(hipFree(Hmm)); + REQUIRE(IfTestPassed); + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + +// The following function tests various scenarios around the flag +// 'hipMemAdviseSetReadMostly' using HMM memory and hipMemAdvise() api + +TEST_CASE("Unit_hipMemAdvise_ReadMostly") { + int MangdMem = HmmAttrPrint(); + if (MangdMem == 1) { + bool IfTestPassed = true; + int MEM_SIZE = 4096, A_CONST = 9999; + float *Hmm = nullptr; + HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE)); + for (uint64_t i = 0; i < (MEM_SIZE/sizeof(float)); ++i) { + Hmm[i] = A_CONST; + } + HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetReadMostly, 0)); + // Checking if the data can be read after setting hipMemAdviseSetReadMostly + for (uint64_t i = 0; i < (MEM_SIZE/sizeof(float)); ++i) { + if (Hmm[i] != A_CONST) { + WARN("Didn't find expected value in Hmm memory after setting"); + WARN(" hipMemAdviseSetReadMostly flag line no.: " << __LINE__); + IfTestPassed = false; + } + } + + // Checking if the memory region can be modified + for (uint64_t i = 0; i < (MEM_SIZE/sizeof(float)); ++i) { + Hmm[i] = A_CONST; + } + + for (uint64_t i = 0; i < (MEM_SIZE/sizeof(float)); ++i) { + if (Hmm[i] != A_CONST) { + WARN("Didn't find expected value in Hmm memory after Modification\n"); + WARN("line no.: " << __LINE__); + IfTestPassed = false; + } + } + + int out = A_CONST; + HIP_CHECK(hipMemRangeGetAttribute(&out, 4, hipMemRangeAttributeReadMostly, + Hmm, MEM_SIZE)); + if (out != 1) { + WARN("out value: " << out); + IfTestPassed = false; + } + // Checking the advise attribute after prefetch + HIP_CHECK(hipMemPrefetchAsync(Hmm, MEM_SIZE, 0, 0)); + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipMemRangeGetAttribute(&out, sizeof(int), + hipMemRangeAttributeReadMostly, Hmm, + MEM_SIZE)); + if (out != 1) { + WARN("Attribute assigned to memory changed after calling "); + WARN("hipMemPrefetchAsync(). line no.: " << __LINE__); + WARN("out value: " << out); + IfTestPassed = false; + } + HIP_CHECK(hipFree(Hmm)); + REQUIRE(IfTestPassed); + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + +// The following function verifies if assigning of a flag invalidates the +// earlier flag which was assigned to the same memory region using +// hipMemAdvise() +TEST_CASE("Unit_hipMemAdvise_TstFlgOverrideEffect") { + int MangdMem = HmmAttrPrint(); + if (MangdMem == 1) { + bool IfTestPassed = true; + int MEM_SIZE = 4*4096, A_CONST = 9999; + float *Hmm = nullptr; + int NumDevs = 0, dev = A_CONST; + + HIP_CHECK(hipGetDeviceCount(&NumDevs)); + HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE, hipMemAttachGlobal)); + for (int i = 0; i < NumDevs; ++i) { + HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetReadMostly, i)); + HIP_CHECK(hipMemRangeGetAttribute(&dev, sizeof(int), + hipMemRangeAttributeReadMostly, Hmm, + MEM_SIZE)); + if (dev != 1) { + WARN("hipMemAdviseSetReadMostly flag did not take affect despite "); + WARN("setting it using hipMemAdvise(). line no.: " << __LINE__); + IfTestPassed = false; + break; + } + dev = A_CONST; + HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetPreferredLocation, + i)); + HIP_CHECK(hipMemRangeGetAttribute(&dev, sizeof(int), + hipMemRangeAttributePreferredLocation, + Hmm, MEM_SIZE)); + if (dev != i) { + WARN("hipMemAdviseSetPreferredLocation flag did not take affect "); + WARN("despite setting it using hipMemAdvise()\n"); + WARN("line no.: " << __LINE__); + IfTestPassed = false; + break; + } + + HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetAccessedBy, i)); + dev = A_CONST; + HIP_CHECK(hipMemRangeGetAttribute(&dev, sizeof(int), + hipMemRangeAttributeAccessedBy, Hmm, + MEM_SIZE)); + if (dev != i) { + WARN("hipMemAdviseSetAccessedBy flag did not take affect despite "); + WARN("setting it using hipMemAdvise(). line no.: " << __LINE__); + IfTestPassed = false; + break; + } + HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseUnsetAccessedBy, i)); + } + HIP_CHECK(hipFree(Hmm)); + REQUIRE(IfTestPassed); + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + + +// The following function tests if peers can set hipMemAdviseSetAccessedBy flag +// on HMM memory prefetched on each of the other gpus +#if HT_AMD +TEST_CASE("Unit_hipMemAdvise_TstAccessedByPeer") { + int MangdMem = HmmAttrPrint(); + if (MangdMem == 1) { + bool IfTestPassed = true; + int *Hmm = nullptr, MEM_SIZE = 4*4096, A_CONST = 9999;; + int NumDevs = 0, CanAccessPeer = A_CONST, flag = 0; + + HIP_CHECK(hipGetDeviceCount(&NumDevs)); + if (NumDevs < 2) { + SUCCEED("Test TestSetAccessedByPeer() need atleast two Gpus to test" + " the scenario. This system has GPUs less than 2"); + } + HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE, hipMemAttachGlobal)); + for (int i = 0; i < NumDevs; ++i) { + HIP_CHECK(hipMemPrefetchAsync(Hmm, MEM_SIZE, i, 0)); + for (int j = 0; j < NumDevs; ++j) { + if (i == j) + continue; + HIP_CHECK(hipSetDevice(j)); + HIP_CHECK(hipDeviceCanAccessPeer(&CanAccessPeer, j, i)); + if (CanAccessPeer) { + HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetAccessedBy, j)); + for (uint64_t m = 0; m < (MEM_SIZE/sizeof(int)); ++m) { + Hmm[m] = 4; + } + HIP_CHECK(hipDeviceEnablePeerAccess(i, 0)); + MemAdvseKernel<<<(MEM_SIZE/sizeof(int)/32), 32>>>( + (MEM_SIZE/sizeof(int)), Hmm); + HIP_CHECK(hipDeviceSynchronize()); + // Verifying the result + for (uint64_t m = 0; m < (MEM_SIZE/sizeof(int)); ++m) { + if (Hmm[m] != 16) { + flag = 1; + } + } + if (flag) { + WARN("Didnt get Expected results with device: " << j); + WARN("line no.: " << __LINE__); + IfTestPassed = false; + flag = 0; + } + } + } + } + HIP_CHECK(hipFree(Hmm)); + REQUIRE(IfTestPassed); + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} +#endif + + +/* Set AccessedBy flag and check value returned by hipMemRangeGetAttribute() + It should be -2(same is observed on cuda)*/ +TEST_CASE("Unit_hipMemAdvise_TstAccessedByFlg") { + int managed = HmmAttrPrint(); + if (managed == 1) { + int *Hmm = NULL, data = 999; + HIP_CHECK(hipMallocManaged(&Hmm, 2*4096)); + HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetAccessedBy, 0)); + HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), + hipMemRangeAttributeLastPrefetchLocation, + Hmm, 2*4096)); + if (data != -2) { + WARN("Didnt get expected value!!\n"); + REQUIRE(false); + } + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + +/* Set AccessedBy flag to device 0 on Hmm memory and prefetch the memory to + device 1, then probe for AccessedBy flag using hipMemRangeGetAttribute() + we should still see the said flag is set for device 0*/ +TEST_CASE("Unit_hipMemAdvise_TstAccessedByFlg2") { + int managed = HmmAttrPrint(); + if (managed == 1) { + int *Hmm = NULL, data = 999, Ngpus = 0; + HIP_CHECK(hipGetDeviceCount(&Ngpus)); + if (Ngpus >= 2) { + hipStream_t strm; + HIP_CHECK(hipStreamCreate(&strm)); + HIP_CHECK(hipMallocManaged(&Hmm, 2*4096)); + HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetAccessedBy, 0)); + HIP_CHECK(hipMemPrefetchAsync(Hmm, 2*4096, 1, strm)); + HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), + hipMemRangeAttributeAccessedBy, Hmm, 2*4096)); + if (data != 0) { + WARN("Didnt get expected behavior at line: " << __LINE__); + REQUIRE(false); + } + HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseUnsetAccessedBy, 0)); + HIP_CHECK(hipStreamDestroy(strm)); + HIP_CHECK(hipFree(Hmm)); + } + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + + + +/* 1) Set AccessedBy to device 0 followed by PreferredLocation to device 1 + check for AccessedBy flag using hipMemRangeGetAttribute() it should + return 0 + 2) Unset AccessedBy to 0 and set it to device 1 followed by + PreferredLocation to device 1, check for AccessedBy flag using + hipMemRangeGetAttribute() it should return 1*/ + +TEST_CASE("Unit_hipMemAdvise_TstAccessedByFlg3") { + int managed = HmmAttrPrint(); + if (managed == 1) { + int *Hmm = NULL, data = 999, Ngpus = 0; + HIP_CHECK(hipGetDeviceCount(&Ngpus)); + if (Ngpus >= 2) { + HIP_CHECK(hipMallocManaged(&Hmm, 2*4096)); + HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetAccessedBy, 0)); + HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetPreferredLocation, 1)); + HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), + hipMemRangeAttributeAccessedBy, Hmm, 2*4096)); + if (data != 0) { + WARN("Didnt get expected behavior at line: " << __LINE__); + REQUIRE(false); + } + HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseUnsetAccessedBy, 0)); + HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetAccessedBy, 1)); + HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetPreferredLocation, 0)); + HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), + hipMemRangeAttributeAccessedBy, Hmm, 2*4096)); + if (data != 1) { + WARN("Didnt get expected behavior at line: " << __LINE__); + REQUIRE(false); + } + HIP_CHECK(hipFree(Hmm)); + } + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + + +/* Set AccessedBy flag to HMM memory launch a kernel and then unset + AccessedBy, launch kernel. We should not have any access issues*/ + +TEST_CASE("Unit_hipMemAdvise_TstAccessedByFlg4") { + int managed = HmmAttrPrint(); + if (managed == 1) { + int *Hmm = NULL, NumElms = (1024 * 1024), InitVal = 123, blockSize = 64; + int DataMismatch = 0; + hipStream_t strm; + HIP_CHECK(hipStreamCreate(&strm)); + HIP_CHECK(hipMallocManaged(&Hmm, (NumElms * sizeof(int)))); + HIP_CHECK(hipMemAdvise(Hmm, (NumElms * sizeof(int)), + hipMemAdviseSetAccessedBy, 0)); + // Initializing memory + for (int i = 0; i < NumElms; ++i) { + Hmm[i] = InitVal; + } + dim3 dimBlock(blockSize, 1, 1); + dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1); + // launching kernel from each one of the gpus + MemAdvise2<<>>(Hmm, NumElms); + HIP_CHECK(hipStreamSynchronize(strm)); + + // verifying the final result + for (int i = 0; i < NumElms; ++i) { + if (Hmm[i] != (InitVal + 10)) { + DataMismatch++; + } + } + + if (DataMismatch != 0) { + WARN("DataMismatch is observed at line: " << __LINE__); + REQUIRE(false); + } + + HIP_CHECK(hipMemAdvise(Hmm, (NumElms * sizeof(int)), + hipMemAdviseUnsetAccessedBy, 0)); + MemAdvise2<<>>(Hmm, NumElms); + HIP_CHECK(hipStreamSynchronize(strm)); + // verifying the final result + for (int i = 0; i < NumElms; ++i) { + if (Hmm[i] != (InitVal + (2*10))) { + DataMismatch++; + } + } + + if (DataMismatch != 0) { + WARN("DataMismatch is observed at line: " << __LINE__); + REQUIRE(false); + } + HIP_CHECK(hipFree(Hmm)); + HIP_CHECK(hipStreamDestroy(strm)); + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + + +/* Allocate memory using aligned_alloc(), assign PreferredLocation flag to + the allocated memory and launch a kernel. Kernel should get executed + successfully without hang or segfault*/ +#if __linux__ && HT_AMD +TEST_CASE("Unit_hipMemAdvise_TstAlignedAllocMem") { + if ((setenv("HSA_XNACK", "1", 1)) != 0) { + WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!"); + REQUIRE(false); + } + // The following code block checks for gfx90a so as to skip if the device is not MI200 + + hipDeviceProp_t prop; + int device; + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipGetDeviceProperties(&prop, device)); + std::string gfxName(prop.gcnArchName); + + if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) { + int stat = 0; + if (fork() == 0) { + // The below part should be inside fork + int managedMem = 0, pageMemAccess = 0; + HIP_CHECK(hipDeviceGetAttribute(&pageMemAccess, + hipDeviceAttributePageableMemoryAccess, 0)); + WARN("hipDeviceAttributePageableMemoryAccess:" << pageMemAccess); + + HIP_CHECK(hipDeviceGetAttribute(&managedMem, hipDeviceAttributeManagedMemory, 0)); + WARN("hipDeviceAttributeManagedMemory: " << managedMem); + if ((managedMem == 1) && (pageMemAccess == 1)) { + int *Mllc = nullptr, MemSz = 4096 * 4, NumElms = 4096, InitVal = 123; + // Mllc = reinterpret_cast<(int *)>(aligned_alloc(4096, MemSz)); + Mllc = reinterpret_cast(aligned_alloc(4096, 4096*4)); + for (int i = 0; i < NumElms; ++i) { + Mllc[i] = InitVal; + } + hipStream_t strm; + int DataMismatch = 0; + HIP_CHECK(hipStreamCreate(&strm)); + // The following hipMemAdvise() call is made to know if advise on + // aligned_alloc() is causing any issue + HIP_CHECK(hipMemAdvise(Mllc, MemSz, hipMemAdviseSetPreferredLocation, 0)); + HIP_CHECK(hipMemPrefetchAsync(Mllc, MemSz, 0, strm)); + HIP_CHECK(hipStreamSynchronize(strm)); + MemAdvise2<<<(NumElms/32), 32, 0, strm>>>(Mllc, NumElms); + HIP_CHECK(hipStreamSynchronize(strm)); + for (int i = 0; i < NumElms; ++i) { + if (Mllc[i] != (InitVal + 10)) { + DataMismatch++; + } + } + if (DataMismatch != 0) { + WARN("DataMismatch observed!!"); + exit(9); // 9 for failure + } else { + exit(10); // 10 for Pass result + } + } else { + SUCCEED("GPU 0 doesn't support ManagedMemory with hipDeviceAttributePageableMemoryAccess " + "attribute. Hence skipping the testing with Pass result.\n"); + exit(Catch::ResultDisposition::ContinueOnFailure); + } + } else { + wait(&stat); + int Result = WEXITSTATUS(stat); + if (Result == Catch::ResultDisposition::ContinueOnFailure) { + WARN("GPU 0 doesn't support ManagedMemory with hipDeviceAttributePageableMemoryAccess " + "attribute. Hence skipping the testing with Pass result.\n"); + } else { + if (Result != 10) { + REQUIRE(false); + } + } + } + } else { + SUCCEED("Memory model feature is only supported for gfx90a, Hence" + "skipping the testcase for this GPU " << device); + WARN("Memory model feature is only supported for gfx90a, Hence" + "skipping the testcase for this GPU " << device); + } + +} +#endif + +/* Allocate Hmm memory, set advise to PreferredLocation and then get + attribute using the api hipMemRangeGetAttribute() for + hipMemRangeAttributeLastPrefetchLocation the value returned should be -2*/ + +TEST_CASE("Unit_hipMemAdvise_TstMemAdvisePrefrdLoc") { + int managed = HmmAttrPrint(); + if (managed == 1) { + int *Hmm = NULL, data = 999; + HIP_CHECK(hipMallocManaged(&Hmm, 4096)); + HIP_CHECK(hipMemAdvise(Hmm, 4096, hipMemAdviseSetPreferredLocation, 0)); + HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), + hipMemRangeAttributeLastPrefetchLocation, + Hmm, 4096)); + if (data != -2) { + WARN("Didnt receive expected value."); + REQUIRE(false); + } + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + + +/* Allocate HMM memory, set PreferredLocation to device 0, Prfetch the mem + to device1, probe for hipMemRangeAttributeLastPrefetchLocation using + hipMemRangeGetAttribute(), we should get 1*/ + +TEST_CASE("Unit_hipMemAdvise_TstMemAdviseLstPreftchLoc") { + int NumDevs = 0; + HIP_CHECK(hipGetDeviceCount(&NumDevs)); + if (NumDevs >= 2) { + int managed = HmmAttrPrint(); + if (managed == 1) { + int *Hmm = NULL, data = 999; + hipStream_t strm; + HIP_CHECK(hipSetDevice(1)); + HIP_CHECK(hipStreamCreate(&strm)); + HIP_CHECK(hipMallocManaged(&Hmm, 4096)); + HIP_CHECK(hipMemAdvise(Hmm, 4096, hipMemAdviseSetPreferredLocation, 0)); + HIP_CHECK(hipMemPrefetchAsync(Hmm, 4096, 1, strm)); + HIP_CHECK(hipStreamSynchronize(strm)); + HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), + hipMemRangeAttributeLastPrefetchLocation, + Hmm, 4096)); + if (data != 1) { + WARN("Didnt receive expected value!!"); + REQUIRE(false); + } + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } + } else { + SUCCEED("This system has less than 2 gpus hence skipping the test.\n"); + } +} + + +/* Allocate HMM memory, set ReadMostly followed by PreferredLocation, probe + for hipMemRangeAttributeReadMostly and hipMemRangeAttributePreferredLocation + using hipMemRangeGetAttribute() we should observe 1 and 0 correspondingly. + In other words setting of hipMemRangeAttributePreferredLocation should not + impact hipMemRangeAttributeReadMostly advise to the memory*/ + +TEST_CASE("Unit_hipMemAdvise_TstMemAdviseMultiFlag") { + int managed = HmmAttrPrint(); + if (managed == 1) { + int *Hmm = NULL, data = 999; + HIP_CHECK(hipMallocManaged(&Hmm, 4096)); + HIP_CHECK(hipMemAdvise(Hmm, 4096, hipMemAdviseSetReadMostly, 0)); + HIP_CHECK(hipMemAdvise(Hmm, 4096, hipMemAdviseSetPreferredLocation, 0)); + HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), + hipMemRangeAttributeReadMostly, Hmm, + 4096)); + if (data != 1) { + WARN("Didnt receive expected value at line: " << data); + REQUIRE(false); + } + HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), + hipMemRangeAttributePreferredLocation, Hmm, + 4096)); + if (data != 0) { + WARN("Didnt receive expected value at line: " << data); + REQUIRE(false); + } + HIP_CHECK(hipFree(Hmm)); + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + + + +/*Allocate Hmm memory, advise it to ReadMostly for gpu: 0 and launch kernel + on all other gpus except 0. This test case may discover any effect or + access denial case arising due to setting ReadMostly only to a particular + gpu*/ + +TEST_CASE("Unit_hipMemAdvise_ReadMosltyMgpuTst") { + int managed = HmmAttrPrint(); + if (managed == 1) { + int Ngpus = 0; + HIP_CHECK(hipGetDeviceCount(&Ngpus)); + if (Ngpus < 2) { + SUCCEED("This test needs atleast two gpus to run." + "Hence skipping the test.\n"); + } + int *Hmm = NULL, NumElms = (1024 * 1024), InitVal = 123, blockSize = 64; + int *Hmm1 = NULL, DataMismatch = 0; + hipStream_t strm; + HIP_CHECK(hipStreamCreate(&strm)); + HIP_CHECK(hipMallocManaged(&Hmm, (NumElms * sizeof(int)))); + // Initializing memory + for (int i = 0; i < NumElms; ++i) { + Hmm[i] = InitVal; + } + HIP_CHECK(hipMemAdvise(Hmm, (NumElms * sizeof(int)), + hipMemAdviseSetReadMostly, 0)); + dim3 dimBlock(blockSize, 1, 1); + dim3 dimGrid((NumElms + blockSize -1)/blockSize, 1, 1); +#if HT_AMD + SECTION("Launch Kernel on all other gpus") { + // launching kernel from each one of the gpus + for (int i = 1; i < Ngpus; ++i) { + DataMismatch = 0; + HIP_CHECK(hipSetDevice(i)); + HIP_CHECK(hipMallocManaged(&Hmm1, (NumElms * sizeof(int)))); + MemAdvise3<<>>(Hmm, Hmm1, NumElms); + HIP_CHECK(hipStreamSynchronize(strm)); + // verifying the results + for (int j = 0; j < NumElms; ++j) { + if (Hmm1[j] != (InitVal + 10)) { + DataMismatch++; + } + } + if (DataMismatch != 0) { + WARN("DataMismatch is observed with the gpu: " << i); + REQUIRE(false); + } + HIP_CHECK(hipFree(Hmm1)); + } + } + + SECTION("Launch Kernel on all other gpus and manipulate the content") { + for (int i = 0; i < Ngpus; ++i) { + DataMismatch = 0; + HIP_CHECK(hipSetDevice(i)); + HIP_CHECK(hipMemAdvise(Hmm, (NumElms * sizeof(int)), + hipMemAdviseSetReadMostly, i)); + MemAdvise2<<>>(Hmm, NumElms); + HIP_CHECK(hipStreamSynchronize(strm)); + } + // verifying the final result + for (int i = 0; i < NumElms; ++i) { + if (Hmm[i] != (InitVal + Ngpus * 10)) { + DataMismatch++; + } + } + + if (DataMismatch != 0) { + WARN("DataMismatch is observed at line: " << __LINE__); + REQUIRE(false); + } + } +#endif + HIP_CHECK(hipFree(Hmm)); + HIP_CHECK(hipStreamDestroy(strm)); + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + + +TEST_CASE("Unit_hipMemAdvise_TstSetUnsetPrfrdLoc") { + int managed = HmmAttrPrint(); + if (managed == 1) { + int *Hmm = NULL, data = 999; + HIP_CHECK(hipMallocManaged(&Hmm, 2*4096)); + HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseSetPreferredLocation, 0)); + HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), + hipMemRangeAttributePreferredLocation, Hmm, 2*4096)); + if (data != 0) { + WARN("Didnt receive expected value!!"); + REQUIRE(false); + } + HIP_CHECK(hipMemAdvise(Hmm, 2*4096, hipMemAdviseUnsetPreferredLocation, 0)); + HIP_CHECK(hipMemRangeGetAttribute(&data, sizeof(int), + hipMemRangeAttributePreferredLocation, Hmm, 2*4096)); + if (data != -2) { + WARN("Didnt receive expected value!!"); + REQUIRE(false); + } + } else { + SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory " + "attribute. Hence skipping the testing with Pass result.\n"); + } +} + From c8ccf14ec97eb504788198f237c97460f8803385 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Wed, 1 Feb 2023 13:22:14 +0530 Subject: [PATCH 03/13] Update CI stage names to include backend label (#150) --- .jenkins/jenkinsfile | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/.jenkins/jenkinsfile b/.jenkins/jenkinsfile index ae4984e5a6..bed5be315e 100644 --- a/.jenkins/jenkinsfile +++ b/.jenkins/jenkinsfile @@ -1,6 +1,6 @@ def hipBuildTest(String backendLabel) { node(backendLabel) { - stage("Source sync ${backendLabel}") { + stage("SYNC - ${backendLabel}") { // Checkout hip repository with the PR patch dir("${WORKSPACE}/hip-tests") { @@ -37,7 +37,7 @@ def hipBuildTest(String backendLabel) { } } - stage("Build - Catch2 framework") { + stage("BUILD HIP - ${backendLabel}") { // Running the build on hipamd workspace dir("${WORKSPACE}/hipamd") { sh """#!/usr/bin/env bash @@ -56,7 +56,7 @@ def hipBuildTest(String backendLabel) { } } - stage("Build - HIP TESTS") { + stage("BUILD HIP TESTS - ${backendLabel}") { // Running the build on HIP TESTS workspace dir("${WORKSPACE}/hip-tests") { env.HIP_PATH = "${HIPAMD_DIR}" + "/build/install" @@ -78,7 +78,7 @@ def hipBuildTest(String backendLabel) { } } - stage('HIP Unit Tests - Catch2 framework') { + stage("TEST - ${backendLabel}") { dir("${WORKSPACE}/hip-tests") { sh """#!/usr/bin/env bash set -x From 9522e9343daf9ffb46eae465ded2447a98450736 Mon Sep 17 00:00:00 2001 From: music-dino <111048524+music-dino@users.noreply.github.com> Date: Wed, 1 Feb 2023 15:11:42 +0100 Subject: [PATCH 04/13] EXSWHTEC-150 - Implement tests for hipImportExternalMemory and hipDestroyExternalMemory for the Vulkan API (#26) - Basic positive tests - Negative parameter tests --- catch/unit/vulkan_interop/CMakeLists.txt | 2 + .../hipDestroyExternalMemory.cc | 45 +++++++++++ .../vulkan_interop/hipImportExternalMemory.cc | 76 +++++++++++++++++++ 3 files changed, 123 insertions(+) create mode 100644 catch/unit/vulkan_interop/hipDestroyExternalMemory.cc create mode 100644 catch/unit/vulkan_interop/hipImportExternalMemory.cc diff --git a/catch/unit/vulkan_interop/CMakeLists.txt b/catch/unit/vulkan_interop/CMakeLists.txt index 2f015d726a..a70940cf5c 100644 --- a/catch/unit/vulkan_interop/CMakeLists.txt +++ b/catch/unit/vulkan_interop/CMakeLists.txt @@ -1,6 +1,8 @@ set(TEST_SRC vulkan_test.cc hipExternalMemoryGetMappedBuffer.cc + hipImportExternalMemory.cc + hipDestroyExternalMemory.cc ) find_package(Vulkan) diff --git a/catch/unit/vulkan_interop/hipDestroyExternalMemory.cc b/catch/unit/vulkan_interop/hipDestroyExternalMemory.cc new file mode 100644 index 0000000000..3f7c3cc13d --- /dev/null +++ b/catch/unit/vulkan_interop/hipDestroyExternalMemory.cc @@ -0,0 +1,45 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "vulkan_test.hh" + +constexpr bool enable_validation = false; + +TEST_CASE("Unit_hipDestroyExternalMemory_Vulkan_Negative_Parameters") { + SECTION("extMem == nullptr") { + HIP_CHECK_ERROR(hipDestroyExternalMemory(nullptr), hipErrorInvalidValue); + } + +// Segfaults in CUDA +// Disabled on AMD due to defect - EXSWHTEC-187 +#if HT_AMD && 0 + SECTION("Double free") { + VulkanTest vkt(enable_validation); + const auto storage = vkt.CreateMappedStorage(1, VK_BUFFER_USAGE_TRANSFER_DST_BIT, true); + auto desc = vkt.BuildMemoryDescriptor(storage.memory, sizeof(*storage.host_ptr)); + hipExternalMemory_t ext_memory; + HIP_CHECK(hipImportExternalMemory(&ext_memory, &desc)); + + HIP_CHECK(hipDestroyExternalMemory(ext_memory)); + HIP_CHECK_ERROR(hipDestroyExternalMemory(ext_memory), hipErrorInvalidValue); + } +#endif +} \ No newline at end of file diff --git a/catch/unit/vulkan_interop/hipImportExternalMemory.cc b/catch/unit/vulkan_interop/hipImportExternalMemory.cc new file mode 100644 index 0000000000..5abfada9d5 --- /dev/null +++ b/catch/unit/vulkan_interop/hipImportExternalMemory.cc @@ -0,0 +1,76 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "vulkan_test.hh" + +constexpr bool enable_validation = false; + +TEST_CASE("Unit_hipImportExternalMemory_Vulkan_Negative_Parameters") { + VulkanTest vkt(enable_validation); + const auto storage = vkt.CreateMappedStorage(1, VK_BUFFER_USAGE_TRANSFER_DST_BIT, true); + auto desc = vkt.BuildMemoryDescriptor(storage.memory, sizeof(*storage.host_ptr)); + hipExternalMemory_t ext_memory; + +// Disabled due to defect - EXSWHTEC-182 +#if HT_NVIDIA + SECTION("extMem_out == nullptr") { + HIP_CHECK_ERROR(hipImportExternalMemory(nullptr, &desc), hipErrorInvalidValue); + } +#endif + +// Disabled due to defect - EXSWHTEC-183 +#if HT_NVIDIA + SECTION("memHandleDesc == nullptr") { + HIP_CHECK_ERROR(hipImportExternalMemory(&ext_memory, nullptr), hipErrorInvalidValue); + } +#endif + +// Disabled due to defect - EXSWHTEC-185 +#if HT_NVIDIA + SECTION("memHandleDesc.size == 0") { + desc.size = 0; + HIP_CHECK_ERROR(hipImportExternalMemory(&ext_memory, &desc), hipErrorInvalidValue); + } +#endif + +// Disabled due to defect - EXSWHTEC-186 +#if HT_NVIDIA + SECTION("Invalid memHandleDesc.flags") { + desc.flags = 2; + HIP_CHECK_ERROR(hipImportExternalMemory(&ext_memory, &desc), hipErrorInvalidValue); + } +#endif + +// Disabled due to defect - EXSWHTEC-184 +#if HT_NVIDIA + SECTION("Invalid memHandleDesc.type") { + desc.type = static_cast(-1); + HIP_CHECK_ERROR(hipImportExternalMemory(&ext_memory, &desc), hipErrorInvalidValue); + } +#endif + +#ifdef _WIN32 + SECTION("memHandleDesc.handle == NULL") { + desc.handle.win32.handle = NULL; + HIP_CHECK_ERROR(hipImportExternalMemory(&ext_memory, &desc), hipErrorInvalidValue); + } +#endif +} \ No newline at end of file From d6f043fe372c5b75c12d289c6df223c0c97e59d5 Mon Sep 17 00:00:00 2001 From: music-dino <111048524+music-dino@users.noreply.github.com> Date: Wed, 1 Feb 2023 15:13:20 +0100 Subject: [PATCH 05/13] EXSWHTEC-147 - Implement tests for hipWaitExternalSemaphoresAsync for the Vulkan API (#28) - Basic positive tests - Negative parameter tests --- catch/unit/vulkan_interop/CMakeLists.txt | 1 + .../hipWaitExternalSemaphoresAsync.cc | 231 ++++++++++++++++++ 2 files changed, 232 insertions(+) create mode 100644 catch/unit/vulkan_interop/hipWaitExternalSemaphoresAsync.cc diff --git a/catch/unit/vulkan_interop/CMakeLists.txt b/catch/unit/vulkan_interop/CMakeLists.txt index a70940cf5c..3421215a83 100644 --- a/catch/unit/vulkan_interop/CMakeLists.txt +++ b/catch/unit/vulkan_interop/CMakeLists.txt @@ -3,6 +3,7 @@ set(TEST_SRC hipExternalMemoryGetMappedBuffer.cc hipImportExternalMemory.cc hipDestroyExternalMemory.cc + hipWaitExternalSemaphoresAsync.cc ) find_package(Vulkan) diff --git a/catch/unit/vulkan_interop/hipWaitExternalSemaphoresAsync.cc b/catch/unit/vulkan_interop/hipWaitExternalSemaphoresAsync.cc new file mode 100644 index 0000000000..edebebe52a --- /dev/null +++ b/catch/unit/vulkan_interop/hipWaitExternalSemaphoresAsync.cc @@ -0,0 +1,231 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "vulkan_test.hh" + +constexpr bool enable_validation = false; + +TEST_CASE("Unit_hipWaitExternalSemaphoresAsync_Vulkan_Positive_Binary_Semaphore") { + VulkanTest vkt(enable_validation); + + constexpr uint32_t count = 1; + const auto src_storage = vkt.CreateMappedStorage(count, VK_BUFFER_USAGE_TRANSFER_SRC_BIT); + const auto dst_storage = vkt.CreateMappedStorage(count, VK_BUFFER_USAGE_TRANSFER_DST_BIT); + + const auto command_buffer = vkt.GetCommandBuffer(); + + VkCommandBufferBeginInfo begin_info = {}; + begin_info.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO; + begin_info.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT; + VK_CHECK_RESULT(vkBeginCommandBuffer(command_buffer, &begin_info)); + VkBufferCopy buffer_copy = {}; + buffer_copy.size = count * sizeof(*src_storage.host_ptr); + vkCmdCopyBuffer(command_buffer, src_storage.buffer, dst_storage.buffer, 1, &buffer_copy); + VK_CHECK_RESULT(vkEndCommandBuffer(command_buffer)); + + const auto semaphore = vkt.CreateExternalSemaphore(VK_SEMAPHORE_TYPE_BINARY); + const auto hip_sem_handle_desc = + vkt.BuildSemaphoreDescriptor(semaphore, VK_SEMAPHORE_TYPE_BINARY); + + hipExternalSemaphore_t hip_ext_semaphore; + HIP_CHECK(hipImportExternalSemaphore(&hip_ext_semaphore, &hip_sem_handle_desc)); + + hipExternalSemaphoreWaitParams hip_ext_semaphore_wait_params = {}; + hip_ext_semaphore_wait_params.flags = 0; + hip_ext_semaphore_wait_params.params.fence.value = 0; + HIP_CHECK(hipWaitExternalSemaphoresAsync(&hip_ext_semaphore, &hip_ext_semaphore_wait_params, 1, + nullptr)); + PollStream(nullptr, hipErrorNotReady); + + VkSubmitInfo submit_info = {}; + submit_info.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO; + submit_info.commandBufferCount = 1; + submit_info.pCommandBuffers = &command_buffer; + submit_info.signalSemaphoreCount = 1; + submit_info.pSignalSemaphores = &semaphore; + + *src_storage.host_ptr = 42; + + const auto fence = vkt.CreateFence(); + VK_CHECK_RESULT(vkQueueSubmit(vkt.GetQueue(), 1, &submit_info, fence)); + VK_CHECK_RESULT( + vkWaitForFences(vkt.GetDevice(), 1, &fence, VK_TRUE, 5'000'000'000 /*5 seconds*/)); + + PollStream(nullptr, hipSuccess); + + REQUIRE(42 == *dst_storage.host_ptr); + + HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore)); +} + +// Timeline semaphores unsupported on AMD +#if HT_NVIDIA +TEST_CASE("Unit_hipWaitExternalSemaphoresAsync_Vulkan_Positive_Timeline_Semaphore") { + VulkanTest vkt(enable_validation); + + const auto [wait_value, signal_value] = + GENERATE(std::make_pair(2, 2), std::make_pair(2, 3), std::make_pair(3, 2)); + INFO("Wait value: " << wait_value << ", signal value: " << signal_value); + + const auto semaphore = vkt.CreateExternalSemaphore(VK_SEMAPHORE_TYPE_TIMELINE); + const auto hip_sem_handle_desc = + vkt.BuildSemaphoreDescriptor(semaphore, VK_SEMAPHORE_TYPE_TIMELINE); + hipExternalSemaphore_t hip_ext_semaphore; + HIP_CHECK(hipImportExternalSemaphore(&hip_ext_semaphore, &hip_sem_handle_desc)); + + hipExternalSemaphoreWaitParams hip_ext_semaphore_wait_params = {}; + hip_ext_semaphore_wait_params.flags = 0; + hip_ext_semaphore_wait_params.params.fence.value = wait_value; + HIP_CHECK(hipWaitExternalSemaphoresAsync(&hip_ext_semaphore, &hip_ext_semaphore_wait_params, 1, + nullptr)); + PollStream(nullptr, hipErrorNotReady); + + VkSemaphoreSignalInfo signal_info = {}; + signal_info.sType = VK_STRUCTURE_TYPE_SEMAPHORE_SIGNAL_INFO; + signal_info.semaphore = semaphore; + signal_info.value = signal_value; + VK_CHECK_RESULT(vkSignalSemaphore(vkt.GetDevice(), &signal_info)); + if (wait_value > signal_value) { + PollStream(nullptr, hipErrorNotReady); + signal_info.value = wait_value; + VK_CHECK_RESULT(vkSignalSemaphore(vkt.GetDevice(), &signal_info)); + } + PollStream(nullptr, hipSuccess); + + HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore)); +} +#endif + +TEST_CASE("Unit_hipWaitExternalSemaphoresAsync_Vulkan_Positive_Multiple_Semaphores") { + VulkanTest vkt(enable_validation); + +#if HT_AMD + constexpr auto second_semaphore_type = VK_SEMAPHORE_TYPE_BINARY; +#else + constexpr auto second_semaphore_type = VK_SEMAPHORE_TYPE_TIMELINE; +#endif + + constexpr uint32_t count = 1; + const auto src_storage = vkt.CreateMappedStorage(count, VK_BUFFER_USAGE_TRANSFER_SRC_BIT); + const auto dst_storage = vkt.CreateMappedStorage(count, VK_BUFFER_USAGE_TRANSFER_DST_BIT); + + const auto command_buffer = vkt.GetCommandBuffer(); + + VkCommandBufferBeginInfo begin_info = {}; + begin_info.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO; + begin_info.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT; + VK_CHECK_RESULT(vkBeginCommandBuffer(command_buffer, &begin_info)); + VkBufferCopy buffer_copy = {}; + buffer_copy.size = count * sizeof(*src_storage.host_ptr); + vkCmdCopyBuffer(command_buffer, src_storage.buffer, dst_storage.buffer, 1, &buffer_copy); + VK_CHECK_RESULT(vkEndCommandBuffer(command_buffer)); + + const auto binary_semaphore = vkt.CreateExternalSemaphore(VK_SEMAPHORE_TYPE_BINARY); + const auto hip_binary_sem_handle_desc = + vkt.BuildSemaphoreDescriptor(binary_semaphore, VK_SEMAPHORE_TYPE_BINARY); + hipExternalSemaphore_t hip_binary_ext_semaphore; + HIP_CHECK(hipImportExternalSemaphore(&hip_binary_ext_semaphore, &hip_binary_sem_handle_desc)); + + const auto timeline_semaphore = vkt.CreateExternalSemaphore(second_semaphore_type); + const auto hip_timeline_sem_handle_desc = + vkt.BuildSemaphoreDescriptor(timeline_semaphore, second_semaphore_type); + hipExternalSemaphore_t hip_timeline_ext_semaphore; + HIP_CHECK(hipImportExternalSemaphore(&hip_timeline_ext_semaphore, &hip_timeline_sem_handle_desc)); + + hipExternalSemaphoreWaitParams binary_semaphore_wait_params = {}; + binary_semaphore_wait_params.params.fence.value = 0; + + hipExternalSemaphoreWaitParams timeline_semaphore_wait_params = {}; + timeline_semaphore_wait_params.params.fence.value = + second_semaphore_type == VK_SEMAPHORE_TYPE_TIMELINE ? 1 : 0; + + hipExternalSemaphore_t ext_semaphores[] = {hip_binary_ext_semaphore, hip_timeline_ext_semaphore}; + hipExternalSemaphoreWaitParams wait_params[] = {binary_semaphore_wait_params, + timeline_semaphore_wait_params}; + HIP_CHECK(hipWaitExternalSemaphoresAsync(ext_semaphores, wait_params, 2, nullptr)); + + PollStream(nullptr, hipErrorNotReady); + + if (second_semaphore_type == VK_SEMAPHORE_TYPE_TIMELINE) { + VkSemaphoreSignalInfo signal_info = {}; + signal_info.sType = VK_STRUCTURE_TYPE_SEMAPHORE_SIGNAL_INFO; + signal_info.semaphore = timeline_semaphore; + signal_info.value = 1; + VK_CHECK_RESULT(vkSignalSemaphore(vkt.GetDevice(), &signal_info)); + + PollStream(nullptr, hipErrorNotReady); + } + + VkSubmitInfo submit_info = {}; + VkSemaphore signal_semaphores[] = {binary_semaphore, timeline_semaphore}; + submit_info.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO; + submit_info.commandBufferCount = 1; + submit_info.pCommandBuffers = &command_buffer; + submit_info.signalSemaphoreCount = second_semaphore_type == VK_SEMAPHORE_TYPE_TIMELINE ? 1 : 2; + submit_info.pSignalSemaphores = + second_semaphore_type == VK_SEMAPHORE_TYPE_MAX_ENUM ? &binary_semaphore : signal_semaphores; + + const auto fence = vkt.CreateFence(); + VK_CHECK_RESULT(vkQueueSubmit(vkt.GetQueue(), 1, &submit_info, fence)); + VK_CHECK_RESULT( + vkWaitForFences(vkt.GetDevice(), 1, &fence, VK_TRUE, 5'000'000'000 /*5 seconds*/)); + + PollStream(nullptr, hipSuccess); + + HIP_CHECK(hipDestroyExternalSemaphore(hip_timeline_ext_semaphore)); + HIP_CHECK(hipDestroyExternalSemaphore(hip_binary_ext_semaphore)); +} + +TEST_CASE("Unit_hipWaitExternalSemaphoresAsync_Vulkan_Negative_Parameters") { + VulkanTest vkt(enable_validation); + hipExternalSemaphoreWaitParams wait_params = {}; + wait_params.params.fence.value = 1; + + SECTION("extSemArray == nullptr") { + HIP_CHECK_ERROR(hipWaitExternalSemaphoresAsync(nullptr, &wait_params, 1, nullptr), + hipErrorInvalidValue); + } + + SECTION("paramsArray == nullptr") { + const auto hip_ext_semaphore = ImportBinarySemaphore(vkt); + HIP_CHECK_ERROR(hipWaitExternalSemaphoresAsync(&hip_ext_semaphore, nullptr, 1, nullptr), + hipErrorInvalidValue); + HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore)); + } + + SECTION("Wait params flag != 0") { + const auto hip_ext_semaphore = ImportBinarySemaphore(vkt); + wait_params.flags = 1; + HIP_CHECK_ERROR(hipWaitExternalSemaphoresAsync(&hip_ext_semaphore, &wait_params, 1, nullptr), + hipErrorInvalidValue); + HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore)); + } + + SECTION("Invalid stream") { + const auto hip_ext_semaphore = ImportBinarySemaphore(vkt); + hipStream_t stream = nullptr; + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK_ERROR(hipWaitExternalSemaphoresAsync(&hip_ext_semaphore, &wait_params, 1, stream), + hipErrorInvalidValue); + HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore)); + } +} \ No newline at end of file From 311380d6bad62b43856e20aa145d76ec2084f1a7 Mon Sep 17 00:00:00 2001 From: music-dino <111048524+music-dino@users.noreply.github.com> Date: Wed, 1 Feb 2023 16:43:51 +0100 Subject: [PATCH 06/13] EXSWHTEC-146 - Implement tests for hipSignalExternalSemaphoresAsync for the Vulkan API (#27) - Basic positive tests - Negative parameter tests --- catch/unit/vulkan_interop/CMakeLists.txt | 1 + .../hipSignalExternalSemaphoresAsync.cc | 204 ++++++++++++++++++ 2 files changed, 205 insertions(+) create mode 100644 catch/unit/vulkan_interop/hipSignalExternalSemaphoresAsync.cc diff --git a/catch/unit/vulkan_interop/CMakeLists.txt b/catch/unit/vulkan_interop/CMakeLists.txt index 3421215a83..8ee4556f23 100644 --- a/catch/unit/vulkan_interop/CMakeLists.txt +++ b/catch/unit/vulkan_interop/CMakeLists.txt @@ -4,6 +4,7 @@ set(TEST_SRC hipImportExternalMemory.cc hipDestroyExternalMemory.cc hipWaitExternalSemaphoresAsync.cc + hipSignalExternalSemaphoresAsync.cc ) find_package(Vulkan) diff --git a/catch/unit/vulkan_interop/hipSignalExternalSemaphoresAsync.cc b/catch/unit/vulkan_interop/hipSignalExternalSemaphoresAsync.cc new file mode 100644 index 0000000000..2c5d055e66 --- /dev/null +++ b/catch/unit/vulkan_interop/hipSignalExternalSemaphoresAsync.cc @@ -0,0 +1,204 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "vulkan_test.hh" + +constexpr bool enable_validation = false; + +TEST_CASE("Unit_hipSignalExternalSemaphoresAsync_Vulkan_Positive_Binary_Semaphore") { + VulkanTest vkt(enable_validation); + + constexpr uint32_t count = 1; + const auto src_storage = vkt.CreateMappedStorage(count, VK_BUFFER_USAGE_TRANSFER_SRC_BIT); + const auto dst_storage = vkt.CreateMappedStorage(count, VK_BUFFER_USAGE_TRANSFER_DST_BIT); + + const auto command_buffer = vkt.GetCommandBuffer(); + VkCommandBufferBeginInfo begin_info = {}; + begin_info.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO; + begin_info.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT; + VK_CHECK_RESULT(vkBeginCommandBuffer(command_buffer, &begin_info)); + VkBufferCopy buffer_copy = {}; + buffer_copy.size = count * sizeof(*src_storage.host_ptr); + vkCmdCopyBuffer(command_buffer, src_storage.buffer, dst_storage.buffer, 1, &buffer_copy); + VK_CHECK_RESULT(vkEndCommandBuffer(command_buffer)); + + const auto semaphore = vkt.CreateExternalSemaphore(VK_SEMAPHORE_TYPE_BINARY); + const auto hip_sem_handle_desc = + vkt.BuildSemaphoreDescriptor(semaphore, VK_SEMAPHORE_TYPE_BINARY); + hipExternalSemaphore_t hip_ext_semaphore; + HIP_CHECK(hipImportExternalSemaphore(&hip_ext_semaphore, &hip_sem_handle_desc)); + + VkSubmitInfo submit_info = {}; + submit_info.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO; + submit_info.commandBufferCount = 1; + submit_info.pCommandBuffers = &command_buffer; + submit_info.waitSemaphoreCount = 1; + submit_info.pWaitSemaphores = &semaphore; + const auto fence = vkt.CreateFence(); + VK_CHECK_RESULT(vkQueueSubmit(vkt.GetQueue(), 1, &submit_info, fence)); + + REQUIRE(vkGetFenceStatus(vkt.GetDevice(), fence) == VK_NOT_READY); + + hipExternalSemaphoreSignalParams signal_params = {}; + signal_params.params.fence.value = 0; + HIP_CHECK(hipSignalExternalSemaphoresAsync(&hip_ext_semaphore, &signal_params, 1, nullptr)); + PollStream(nullptr, hipSuccess); + + VK_CHECK_RESULT( + vkWaitForFences(vkt.GetDevice(), 1, &fence, VK_TRUE, 5'000'000'000 /*5 seconds*/)); + + HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore)); +} + +// Timeline semaphores unsupported on AMD +#if HT_NVIDIA +TEST_CASE("Unit_hipSignalExternalSemaphoresAsync_Vulkan_Positive_Timeline_Semaphore") { + VulkanTest vkt(enable_validation); + constexpr uint64_t signal_value = 2; + + const auto semaphore = vkt.CreateExternalSemaphore(VK_SEMAPHORE_TYPE_TIMELINE); + const auto hip_sem_handle_desc = + vkt.BuildSemaphoreDescriptor(semaphore, VK_SEMAPHORE_TYPE_TIMELINE); + hipExternalSemaphore_t hip_ext_semaphore; + HIP_CHECK(hipImportExternalSemaphore(&hip_ext_semaphore, &hip_sem_handle_desc)); + + hipExternalSemaphoreSignalParams signal_params = {}; + signal_params.params.fence.value = signal_value; + + HIP_CHECK(hipSignalExternalSemaphoresAsync(&hip_ext_semaphore, &signal_params, 1, nullptr)); + PollStream(nullptr, hipSuccess); + + uint64_t sem_value = 0u; + VK_CHECK_RESULT(vkGetSemaphoreCounterValue(vkt.GetDevice(), semaphore, &sem_value)); + + REQUIRE(2 == sem_value); + + HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore)); +} +#endif + +TEST_CASE("Unit_hipSignalExternalSemaphoresAsync_Vulkan_Positive_Multiple_Semaphores") { + VulkanTest vkt(enable_validation); + + constexpr uint32_t count = 1; + const auto src_storage = vkt.CreateMappedStorage(count, VK_BUFFER_USAGE_TRANSFER_SRC_BIT); + const auto dst_storage = vkt.CreateMappedStorage(count, VK_BUFFER_USAGE_TRANSFER_DST_BIT); + +#if HT_AMD + constexpr auto second_semaphore_type = VK_SEMAPHORE_TYPE_BINARY; +#else + constexpr auto second_semaphore_type = VK_SEMAPHORE_TYPE_TIMELINE; +#endif + + const auto command_buffer = vkt.GetCommandBuffer(); + VkCommandBufferBeginInfo begin_info = {}; + begin_info.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO; + begin_info.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT; + VK_CHECK_RESULT(vkBeginCommandBuffer(command_buffer, &begin_info)); + VkBufferCopy buffer_copy = {}; + buffer_copy.size = count * sizeof(*src_storage.host_ptr); + vkCmdCopyBuffer(command_buffer, src_storage.buffer, dst_storage.buffer, 1, &buffer_copy); + VK_CHECK_RESULT(vkEndCommandBuffer(command_buffer)); + + const auto binary_semaphore = vkt.CreateExternalSemaphore(VK_SEMAPHORE_TYPE_BINARY); + const auto hip_binary_sem_handle_desc = + vkt.BuildSemaphoreDescriptor(binary_semaphore, VK_SEMAPHORE_TYPE_BINARY); + hipExternalSemaphore_t hip_binary_ext_semaphore; + HIP_CHECK(hipImportExternalSemaphore(&hip_binary_ext_semaphore, &hip_binary_sem_handle_desc)); + + const auto timeline_semaphore = vkt.CreateExternalSemaphore(second_semaphore_type); + const auto hip_timeline_sem_handle_desc = + vkt.BuildSemaphoreDescriptor(timeline_semaphore, second_semaphore_type); + hipExternalSemaphore_t hip_timeline_ext_semaphore; + HIP_CHECK(hipImportExternalSemaphore(&hip_timeline_ext_semaphore, &hip_timeline_sem_handle_desc)); + + uint64_t wait_values[] = {1, 0}; + VkTimelineSemaphoreSubmitInfo timeline_submit_info = {}; + timeline_submit_info.sType = VK_STRUCTURE_TYPE_TIMELINE_SEMAPHORE_SUBMIT_INFO; + timeline_submit_info.waitSemaphoreValueCount = 2; + timeline_submit_info.pWaitSemaphoreValues = wait_values; + + VkSemaphore wait_semaphores[] = {timeline_semaphore, binary_semaphore}; + VkSubmitInfo submit_info = {}; + submit_info.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO; + submit_info.commandBufferCount = 1; + submit_info.pCommandBuffers = &command_buffer; + submit_info.waitSemaphoreCount = 2; + submit_info.pWaitSemaphores = wait_semaphores; + submit_info.pNext = + second_semaphore_type == VK_SEMAPHORE_TYPE_TIMELINE ? &timeline_submit_info : nullptr; + const auto fence = vkt.CreateFence(); + VK_CHECK_RESULT(vkQueueSubmit(vkt.GetQueue(), 1, &submit_info, fence)); + + REQUIRE(vkGetFenceStatus(vkt.GetDevice(), fence) == VK_NOT_READY); + + hipExternalSemaphoreSignalParams binary_signal_params = {}; + binary_signal_params.params.fence.value = 0; + hipExternalSemaphoreSignalParams timeline_signal_params = {}; + timeline_signal_params.params.fence.value = + second_semaphore_type == VK_SEMAPHORE_TYPE_TIMELINE ? 2 : 0; + hipExternalSemaphore_t ext_semaphores[] = {hip_binary_ext_semaphore, hip_timeline_ext_semaphore}; + hipExternalSemaphoreSignalParams signal_params[] = {binary_signal_params, timeline_signal_params}; + HIP_CHECK(hipSignalExternalSemaphoresAsync(ext_semaphores, signal_params, 2, nullptr)); + + VK_CHECK_RESULT( + vkWaitForFences(vkt.GetDevice(), 1, &fence, VK_TRUE, 5'000'000'000 /*5 seconds*/)); + + HIP_CHECK(hipDestroyExternalSemaphore(hip_binary_ext_semaphore)); + HIP_CHECK(hipDestroyExternalSemaphore(hip_timeline_ext_semaphore)); +} + +TEST_CASE("Unit_hipSignalExternalSemaphoresAsync_Vulkan_Negative_Parameters") { + VulkanTest vkt(enable_validation); + hipExternalSemaphoreSignalParams signal_params = {}; + signal_params.params.fence.value = 1; + + SECTION("extSemArray == nullptr") { + HIP_CHECK_ERROR(hipSignalExternalSemaphoresAsync(nullptr, &signal_params, 1, nullptr), + hipErrorInvalidValue); + } + + SECTION("paramsArray == nullptr") { + const auto hip_ext_semaphore = ImportBinarySemaphore(vkt); + HIP_CHECK_ERROR(hipSignalExternalSemaphoresAsync(&hip_ext_semaphore, nullptr, 1, nullptr), + hipErrorInvalidValue); + HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore)); + } + + SECTION("Wait params flags != 0") { + const auto hip_ext_semaphore = ImportBinarySemaphore(vkt); + signal_params.flags = 1; + HIP_CHECK_ERROR( + hipSignalExternalSemaphoresAsync(&hip_ext_semaphore, &signal_params, 1, nullptr), + hipErrorInvalidValue); + HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore)); + } + + SECTION("Invalid stream") { + const auto hip_ext_semaphore = ImportBinarySemaphore(vkt); + hipStream_t stream = nullptr; + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK_ERROR(hipSignalExternalSemaphoresAsync(&hip_ext_semaphore, &signal_params, 1, stream), + hipErrorInvalidValue); + HIP_CHECK(hipDestroyExternalSemaphore(hip_ext_semaphore)); + } +} \ No newline at end of file From 06726393d4396bea6464090b4f2352a261a65e95 Mon Sep 17 00:00:00 2001 From: music-dino <111048524+music-dino@users.noreply.github.com> Date: Wed, 1 Feb 2023 16:45:05 +0100 Subject: [PATCH 07/13] EXSWHTEC-148 - Implement tests for hipImportExternalSemaphore and hipDestroyExternalSemaphore for the Vulkan API (#29) - Basic positive test - Negative parameter tests --- catch/unit/vulkan_interop/CMakeLists.txt | 2 + .../hipDestroyExternalSemaphore.cc | 40 +++++++++++++ .../hipImportExternalSemaphore.cc | 56 +++++++++++++++++++ 3 files changed, 98 insertions(+) create mode 100644 catch/unit/vulkan_interop/hipDestroyExternalSemaphore.cc create mode 100644 catch/unit/vulkan_interop/hipImportExternalSemaphore.cc diff --git a/catch/unit/vulkan_interop/CMakeLists.txt b/catch/unit/vulkan_interop/CMakeLists.txt index 8ee4556f23..4803737598 100644 --- a/catch/unit/vulkan_interop/CMakeLists.txt +++ b/catch/unit/vulkan_interop/CMakeLists.txt @@ -5,6 +5,8 @@ set(TEST_SRC hipDestroyExternalMemory.cc hipWaitExternalSemaphoresAsync.cc hipSignalExternalSemaphoresAsync.cc + hipImportExternalSemaphore.cc + hipDestroyExternalSemaphore.cc ) find_package(Vulkan) diff --git a/catch/unit/vulkan_interop/hipDestroyExternalSemaphore.cc b/catch/unit/vulkan_interop/hipDestroyExternalSemaphore.cc new file mode 100644 index 0000000000..dcef2d7fad --- /dev/null +++ b/catch/unit/vulkan_interop/hipDestroyExternalSemaphore.cc @@ -0,0 +1,40 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "vulkan_test.hh" + +constexpr bool enable_validation = false; + +TEST_CASE("Unit_hipDestroyExternalSemaphore_Vulkan_Negative_Parameters") { + SECTION("extSem == nullptr") { + HIP_CHECK_ERROR(hipDestroyExternalSemaphore(nullptr), hipErrorInvalidValue); + } + +// Segfaults in CUDA +#if HT_AMD + SECTION("Double free") { + VulkanTest vkt(enable_validation); + const auto ext_semaphore = ImportBinarySemaphore(vkt); + HIP_CHECK(hipDestroyExternalSemaphore(ext_semaphore)); + HIP_CHECK_ERROR(hipDestroyExternalSemaphore(ext_semaphore), hipErrorInvalidValue); + } +#endif +} \ No newline at end of file diff --git a/catch/unit/vulkan_interop/hipImportExternalSemaphore.cc b/catch/unit/vulkan_interop/hipImportExternalSemaphore.cc new file mode 100644 index 0000000000..2920645459 --- /dev/null +++ b/catch/unit/vulkan_interop/hipImportExternalSemaphore.cc @@ -0,0 +1,56 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "vulkan_test.hh" + +constexpr bool enable_validation = false; + +TEST_CASE("Unit_hipImportExternalSemaphore_Vulkan_Negative_Parameters") { + VulkanTest vkt(enable_validation); + const auto semaphore = vkt.CreateExternalSemaphore(VK_SEMAPHORE_TYPE_BINARY); + auto handle_desc = vkt.BuildSemaphoreDescriptor(semaphore, VK_SEMAPHORE_TYPE_BINARY); + hipExternalSemaphore_t ext_semaphore; + + SECTION("extSem_out == nullptr") { + HIP_CHECK_ERROR(hipImportExternalSemaphore(nullptr, &handle_desc), hipErrorInvalidValue); + } + + SECTION("semHandleDesc == nullptr") { + HIP_CHECK_ERROR(hipImportExternalSemaphore(&ext_semaphore, nullptr), hipErrorInvalidValue); + } + + SECTION("semHandleDesc.flags != 0") { + handle_desc.flags = 1; + HIP_CHECK_ERROR(hipImportExternalSemaphore(&ext_semaphore, &handle_desc), hipErrorInvalidValue); + } + + SECTION("Invalid semHandleDesc.type") { + handle_desc.type = static_cast(-1); + HIP_CHECK_ERROR(hipImportExternalSemaphore(&ext_semaphore, &handle_desc), hipErrorInvalidValue); + } + +#ifdef _WIN32 + SECTION("semHandleDesc.handle == NULL") { + handle_desc.handle.win32.handle = NULL; + HIP_CHECK_ERROR(hipImportExternalSemaphore(&ext_semaphore, &handle_desc), hipErrorInvalidValue); + } +#endif +} \ No newline at end of file From 7e331b247995988218261db8469ccf7f3707eb80 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Fri, 3 Feb 2023 20:39:44 +0530 Subject: [PATCH 08/13] SWDEV-371483 - Fixed the build of some tests in HIP-Samples (#120) Change-Id: Iab0c79eafb9c919b65e7cb42ac865ce830a9e48d --- .../12_cmake_hip_add_executable/CMakeLists.txt | 6 +++--- .../2_Cookbook/16_assembly_to_executable/Makefile | 8 ++++---- samples/2_Cookbook/17_llvm_ir_to_executable/Makefile | 12 ++++++------ 3 files changed, 13 insertions(+), 13 deletions(-) diff --git a/samples/2_Cookbook/12_cmake_hip_add_executable/CMakeLists.txt b/samples/2_Cookbook/12_cmake_hip_add_executable/CMakeLists.txt index f1a8bf8a82..248dd355b0 100644 --- a/samples/2_Cookbook/12_cmake_hip_add_executable/CMakeLists.txt +++ b/samples/2_Cookbook/12_cmake_hip_add_executable/CMakeLists.txt @@ -25,12 +25,12 @@ endif () if(NOT DEFINED HIP_PATH) if(NOT DEFINED ENV{HIP_PATH}) - set(HIP_PATH "${ROCM_PATH}/hip" CACHE PATH "Path to which HIP has been installed") + set(HIP_PATH ${ROCM_PATH} CACHE PATH "Path to which HIP has been installed") else() set(HIP_PATH $ENV{HIP_PATH} CACHE PATH "Path to which HIP has been installed") endif() endif() -set(CMAKE_MODULE_PATH "${HIP_PATH}/cmake" ${CMAKE_MODULE_PATH}) +set(CMAKE_MODULE_PATH "${HIP_PATH}/hip/cmake" ${CMAKE_MODULE_PATH}) set(CMAKE_HIP_ARCHITECTURES OFF) project(12_cmake) @@ -53,7 +53,7 @@ set_source_files_properties(${MY_SOURCE_FILES} PROPERTIES HIP_SOURCE_PROPERTY_FO hip_add_executable(${MY_TARGET_NAME} ${MY_SOURCE_FILES} HIPCC_OPTIONS ${MY_HIPCC_OPTIONS} HCC_OPTIONS ${MY_HCC_OPTIONS} CLANG_OPTIONS ${MY_CLANG_OPTIONS} NVCC_OPTIONS ${MY_NVCC_OPTIONS}) # Search for rocm in common locations -list(APPEND CMAKE_PREFIX_PATH ${HIP_PATH} ${ROCM_PATH}) +list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH}/hip ${ROCM_PATH}) find_package(hip QUIET) if(TARGET hip::host) message(STATUS "Found hip::host at ${hip_DIR}") diff --git a/samples/2_Cookbook/16_assembly_to_executable/Makefile b/samples/2_Cookbook/16_assembly_to_executable/Makefile index b82ec8fdbe..0a32aa2b9c 100644 --- a/samples/2_Cookbook/16_assembly_to_executable/Makefile +++ b/samples/2_Cookbook/16_assembly_to_executable/Makefile @@ -21,15 +21,15 @@ ifeq ($(OS),Windows_NT) $(error Makefile is not supported on windows platform. Please use cmake instead to build sample.) endif ROCM_PATH?= $(wildcard /opt/rocm/) -HIP_PATH?= $(wildcard $(ROCM_PATH)/hip) +HIP_PATH?= $(ROCM_PATH) ifeq (,$(HIP_PATH)) HIP_PATH=../../.. endif HIPCC=$(HIP_PATH)/bin/hipcc -CLANG=$(HIP_PATH)/../llvm/bin/clang -LLVM_MC=$(HIP_PATH)/../llvm/bin/llvm-mc -CLANG_OFFLOAD_BUNDLER=$(HIP_PATH)/../llvm/bin/clang-offload-bundler +CLANG=$(HIP_PATH)/llvm/bin/clang +LLVM_MC=$(HIP_PATH)/llvm/bin/llvm-mc +CLANG_OFFLOAD_BUNDLER=$(HIP_PATH)/llvm/bin/clang-offload-bundler SRCS=square.cpp diff --git a/samples/2_Cookbook/17_llvm_ir_to_executable/Makefile b/samples/2_Cookbook/17_llvm_ir_to_executable/Makefile index 835f4b2789..d68304dcac 100644 --- a/samples/2_Cookbook/17_llvm_ir_to_executable/Makefile +++ b/samples/2_Cookbook/17_llvm_ir_to_executable/Makefile @@ -21,17 +21,17 @@ ifeq ($(OS),Windows_NT) $(error Makefile is not supported on windows platform. Please use cmake instead to build sample.) endif ROCM_PATH?= $(wildcard /opt/rocm/) -HIP_PATH?= $(wildcard $(ROCM_PATH)/hip) +HIP_PATH?= $(ROCM_PATH) ifeq (,$(HIP_PATH)) HIP_PATH=../../.. endif HIPCC=$(HIP_PATH)/bin/hipcc -CLANG=$(HIP_PATH)/../llvm/bin/clang -LLVM_MC=$(HIP_PATH)/../llvm/bin/llvm-mc -CLANG_OFFLOAD_BUNDLER=$(HIP_PATH)/../llvm/bin/clang-offload-bundler -LLVM_AS=$(HIP_PATH)/../llvm/bin/llvm-as -LLVM_DIS=$(HIP_PATH)/../llvm/bin/llvm-dis +CLANG=$(HIP_PATH)/llvm/bin/clang +LLVM_MC=$(HIP_PATH)/llvm/bin/llvm-mc +CLANG_OFFLOAD_BUNDLER=$(HIP_PATH)/llvm/bin/clang-offload-bundler +LLVM_AS=$(HIP_PATH)/llvm/bin/llvm-as +LLVM_DIS=$(HIP_PATH)/llvm/bin/llvm-dis SRCS=square.cpp From f10981bbbc2042b6a58f922ff71de8238ffb4f82 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Mon, 6 Feb 2023 02:44:20 +0530 Subject: [PATCH 09/13] SWDEV-371332 - remove check for non device pointers (#145) Change-Id: Ib7e83ab1bca8c61e3710cb9dd75d9f932159476a --- catch/unit/memory/hipPointerGetAttributes.cc | 7 ------- 1 file changed, 7 deletions(-) diff --git a/catch/unit/memory/hipPointerGetAttributes.cc b/catch/unit/memory/hipPointerGetAttributes.cc index d4ba5960ce..cbc4554b3a 100644 --- a/catch/unit/memory/hipPointerGetAttributes.cc +++ b/catch/unit/memory/hipPointerGetAttributes.cc @@ -260,8 +260,6 @@ TEST_CASE("Unit_hipPointerGetAttributes_Basic") { REQUIRE(attribs.devicePointer != attribs2.devicePointer); } HIP_CHECK(hipFree(A_d)); - e = hipPointerGetAttributes(&attribs, A_d); - REQUIRE(e == hipErrorInvalidValue); // Device-visible host memory printf("\nDevice-visible host memory (hipHostMalloc)\n"); @@ -272,15 +270,10 @@ TEST_CASE("Unit_hipPointerGetAttributes_Basic") { char* ptr1 = reinterpret_cast(attribs.hostPointer); REQUIRE((ptr1 + Nbytes / 2) == reinterpret_cast(attribs2.hostPointer)); - HIP_CHECK(hipHostFree(A_Pinned_h)); - e = hipPointerGetAttributes(&attribs, A_Pinned_h); - REQUIRE(e == hipErrorInvalidValue); // OS memory printf("\nOS-allocated memory (malloc)\n"); - e = hipPointerGetAttributes(&attribs, A_OSAlloc_h); - REQUIRE(e == hipErrorInvalidValue); } TEST_CASE("Unit_hipPointerGetAttributes_ClusterAlloc") { From d90c401529d6fc1d0a0912f87c671e765d7d7e3b Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Mon, 6 Feb 2023 07:34:27 +0530 Subject: [PATCH 10/13] SWDEV-366390 - [catch2][dtest] Adding additional functional tests for hipGraphGetRootNodes() (#100) Change-Id: Ie9917780a4ba28639f509142141718f5dcf24ffc --- catch/unit/graph/hipGraphGetRootNodes.cc | 120 ++++++++++++++++++++++- 1 file changed, 119 insertions(+), 1 deletion(-) diff --git a/catch/unit/graph/hipGraphGetRootNodes.cc b/catch/unit/graph/hipGraphGetRootNodes.cc index 2be8ea5110..43b900772f 100644 --- a/catch/unit/graph/hipGraphGetRootNodes.cc +++ b/catch/unit/graph/hipGraphGetRootNodes.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal in the Software without restriction, including without limitation the rights @@ -28,6 +28,11 @@ Functional :: nodes list will be set to NULL, and the number of nodes actually obtained will be returned in NumRootNodes. 4) Create a graph with stream capture done on multiple dependent streams. Verify root nodes of created graph are matching the operations pushed which doesn't have dependencies. + 5) Functional Test to validate number of root nodes when dependencies in the graph are dynamically varied. + 6) Functional Test to validate number of root nodes when dependencies in the graph are dynamically varied + in a cloned graph. + 7) Functional Test to validate number of root nodes when a graph with N independent nodes is added as a + child node to another graph. Argument Validation :: 1) Pass graph as nullptr and verify api returns error code. @@ -41,6 +46,12 @@ Argument Validation :: #include #include +#define NUM_OF_DUMMY_NODES 8 + +static __global__ void dummyKernel() { + return; +} + /** * Functional Test for API fetching root node list */ @@ -346,3 +357,110 @@ TEST_CASE("Unit_hipGraphGetRootNodes_ParamValidation") { HIP_CHECK(hipFree(A_d)); HIP_CHECK(hipFree(C_d)); } + +/** + * Functional Test to validate number of root nodes when dependencies + * in the graph are dynamically varied. + */ +TEST_CASE("Unit_hipGraphGetRootNodes_Complx_NumRootNodes") { + hipGraph_t graph; + hipGraphNode_t kernelnode[NUM_OF_DUMMY_NODES]; + hipKernelNodeParams kernelNodeParams[NUM_OF_DUMMY_NODES]; + HIP_CHECK(hipGraphCreate(&graph, 0)); + // Create graph with no dependencies + for (int i = 0; i < NUM_OF_DUMMY_NODES; i++) { + void* kernelArgs[] = {nullptr}; + kernelNodeParams[i].func = reinterpret_cast(dummyKernel); + kernelNodeParams[i].gridDim = dim3(1); + kernelNodeParams[i].blockDim = dim3(1); + kernelNodeParams[i].sharedMemBytes = 0; + kernelNodeParams[i].kernelParams = reinterpret_cast(kernelArgs); + kernelNodeParams[i].extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernelnode[i], graph, nullptr, + 0, &kernelNodeParams[i])); + } + size_t numRootNodes{}; + HIP_CHECK(hipGraphGetRootNodes(graph, nullptr, &numRootNodes)); + REQUIRE(numRootNodes == NUM_OF_DUMMY_NODES); + // Start creating dependencies in a chain + for (size_t i = 0; i < (NUM_OF_DUMMY_NODES - 1); i++) { + numRootNodes = 0; + HIP_CHECK(hipGraphAddDependencies(graph, &kernelnode[i], + &kernelnode[i+1], 1)); + HIP_CHECK(hipGraphGetRootNodes(graph, nullptr, &numRootNodes)); + REQUIRE(numRootNodes == (NUM_OF_DUMMY_NODES - i - 1)); + } + HIP_CHECK(hipGraphDestroy(graph)); +} + +/** + * Functional Test to validate number of root nodes when dependencies + * in the graph are dynamically varied in a cloned graph. + */ +TEST_CASE("Unit_hipGraphGetRootNodes_Complx_NumRootNodes_ClonedGrph") { + hipGraph_t graph, clonedgraph; + hipGraphNode_t kernelnode[NUM_OF_DUMMY_NODES]; + hipKernelNodeParams kernelNodeParams[NUM_OF_DUMMY_NODES]; + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipGraphCreate(&clonedgraph, 0)); + // Create graph with no dependencies + for (int i = 0; i < NUM_OF_DUMMY_NODES; i++) { + void* kernelArgs[] = {nullptr}; + kernelNodeParams[i].func = reinterpret_cast(dummyKernel); + kernelNodeParams[i].gridDim = dim3(1); + kernelNodeParams[i].blockDim = dim3(1); + kernelNodeParams[i].sharedMemBytes = 0; + kernelNodeParams[i].kernelParams = reinterpret_cast(kernelArgs); + kernelNodeParams[i].extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernelnode[i], graph, nullptr, + 0, &kernelNodeParams[i])); + } + size_t numRootNodes{}; + HIP_CHECK(hipGraphClone(&clonedgraph, graph)); + HIP_CHECK(hipGraphGetRootNodes(clonedgraph, nullptr, &numRootNodes)); + REQUIRE(numRootNodes == NUM_OF_DUMMY_NODES); + // Start creating dependencies in a chain + for (size_t i = 0; i < (NUM_OF_DUMMY_NODES - 1); i++) { + numRootNodes = 0; + hipGraphNode_t node1, node2; + HIP_CHECK(hipGraphNodeFindInClone(&node1, kernelnode[i], clonedgraph)); + HIP_CHECK(hipGraphNodeFindInClone(&node2, kernelnode[i+1], clonedgraph)); + HIP_CHECK(hipGraphAddDependencies(clonedgraph, &node1, &node2, 1)); + HIP_CHECK(hipGraphGetRootNodes(clonedgraph, nullptr, &numRootNodes)); + REQUIRE(numRootNodes == (NUM_OF_DUMMY_NODES - i - 1)); + } + HIP_CHECK(hipGraphDestroy(clonedgraph)); + HIP_CHECK(hipGraphDestroy(graph)); +} + +/** + * Functional Test to validate number of root nodes when a graph with N + * independent nodes is added as a child node to another graph. + */ +TEST_CASE("Unit_hipGraphGetRootNodes_Complx_NRootNodesAsChildGraph") { + hipGraph_t graph, graph1; + hipGraphNode_t kernelnode[NUM_OF_DUMMY_NODES]; + hipKernelNodeParams kernelNodeParams[NUM_OF_DUMMY_NODES]; + hipGraphNode_t child_node; + HIP_CHECK(hipGraphCreate(&graph, 0)); + HIP_CHECK(hipGraphCreate(&graph1, 0)); + // Create graph with no dependencies + for (int i = 0; i < NUM_OF_DUMMY_NODES; i++) { + void* kernelArgs[] = {nullptr}; + kernelNodeParams[i].func = reinterpret_cast(dummyKernel); + kernelNodeParams[i].gridDim = dim3(1); + kernelNodeParams[i].blockDim = dim3(1); + kernelNodeParams[i].sharedMemBytes = 0; + kernelNodeParams[i].kernelParams = reinterpret_cast(kernelArgs); + kernelNodeParams[i].extra = nullptr; + HIP_CHECK(hipGraphAddKernelNode(&kernelnode[i], graph, nullptr, + 0, &kernelNodeParams[i])); + } + HIP_CHECK(hipGraphAddChildGraphNode(&child_node, graph1, + nullptr, 0, graph)); + size_t numRootNodes{}; + HIP_CHECK(hipGraphGetRootNodes(graph1, nullptr, &numRootNodes)); + REQUIRE(numRootNodes == 1); + HIP_CHECK(hipGraphDestroy(graph1)); + HIP_CHECK(hipGraphDestroy(graph)); +} From 5ef26cc476430cd1904f33ec78a485ff0ef7921d Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Tue, 7 Feb 2023 01:37:40 +0530 Subject: [PATCH 11/13] SWDEV-364628 - [catch2][dtest] Test cases added for atomicInc and atomicDec (#121) Change-Id: I2820423db6d588a6b638aeb30b965154843f877c --- catch/unit/deviceLib/Atomic_func.cc | 119 ++++++++++++++++++++++++++++ catch/unit/deviceLib/CMakeLists.txt | 1 + 2 files changed, 120 insertions(+) create mode 100644 catch/unit/deviceLib/Atomic_func.cc diff --git a/catch/unit/deviceLib/Atomic_func.cc b/catch/unit/deviceLib/Atomic_func.cc new file mode 100644 index 0000000000..4a28839c6f --- /dev/null +++ b/catch/unit/deviceLib/Atomic_func.cc @@ -0,0 +1,119 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include +#include +#include + +// Test case to validate atomicInc and atomicDec functions. +// if TestToRun=1, then atomicInc function will be tested and validated +// if TestToRun=2, then atomicDec function will be tested and validated. + + +// kernel function for atomicInc +static __global__ void AtomicCheckInc(int* g_ptr) { + atomicInc(reinterpret_cast(&g_ptr[0]), 17); +} + +// kernel function for atomicDec +static __global__ void AtomicCheckDec(int* g_ptr) { + atomicDec(reinterpret_cast(&g_ptr[0]), 25); +} + +// verify results for atomicInc +static int verifyResultInc(int value) { + int limit = 17; + value = (value >= limit) ? 0 : value + 1; + return value; +} + +// verify results for atomicDec +static int verifyResultDec(int value) { + int limit = 25; + value = ((value == 0) || (value > limit)) ? limit : value - 1; + return value; +} + +// common fuction to launch atomic functions kernel. +static void launchAtomicFunction(int *Hptr, int val, int TestToRun) { + unsigned int memSize = sizeof(int) * 1; + int *dptr{nullptr}; + // allocate device memory + HIP_CHECK(hipMalloc(reinterpret_cast(&dptr), memSize)); + // copy host memory to device + HIP_CHECK(hipMemcpy(dptr, Hptr, memSize, hipMemcpyHostToDevice)); + // launch kernel function + if (TestToRun == 1) { + AtomicCheckInc<<<1, 1>>>(dptr); + } else if (TestToRun == 2) { + AtomicCheckDec<<<1, 1>>>(dptr); + } + // copy back from device to host + HIP_CHECK(hipMemcpy(Hptr, dptr, memSize, hipMemcpyDeviceToHost)); + // verify the results. + if (TestToRun == 1) { + int result = verifyResultInc(val); + REQUIRE(result == Hptr[0]); + } else if (TestToRun == 2) { + int result = verifyResultDec(val); + REQUIRE(result == Hptr[0]); + } + // Cleanup memory + HIP_CHECK(hipFree(dptr)); +} + +TEST_CASE("Unit_AtomicFunctions_Inc") { + int *Hptr{nullptr}; + int val; + // Allocate Host memory + Hptr = reinterpret_cast(malloc(sizeof(int))); + SECTION("Test case when value is lesser than limit") { + val = Hptr[0] = 10; + launchAtomicFunction(Hptr, val, 1); + } + SECTION("Test case when value is greater than limit") { + val = Hptr[0] = 20; + launchAtomicFunction(Hptr, val, 1); + } + SECTION("Test case when value is equal to the limit") { + val = Hptr[0] = 17; + launchAtomicFunction(Hptr, val, 1); + } + free(Hptr); +} + +TEST_CASE("Unit_AtomicFunctions_Dec") { + int *Hptr{nullptr}; + int val; + // Allocate Host memory + Hptr = reinterpret_cast(malloc(sizeof(int))); + SECTION("Test case when value is less than limit") { + val = Hptr[0] = 4; + launchAtomicFunction(Hptr, val, 2); + } + SECTION("Test case when value is greater than limit") { + val = Hptr[0] = 31; + launchAtomicFunction(Hptr, val, 2); + } + SECTION("Test case when value is equal to the limit") { + val = Hptr[0] = 25; + launchAtomicFunction(Hptr, val, 2); + } + free(Hptr); +} diff --git a/catch/unit/deviceLib/CMakeLists.txt b/catch/unit/deviceLib/CMakeLists.txt index 02b542e71e..3ba4bddd39 100644 --- a/catch/unit/deviceLib/CMakeLists.txt +++ b/catch/unit/deviceLib/CMakeLists.txt @@ -13,6 +13,7 @@ set(TEST_SRC syncthreadsand.cc syncthreadscount.cc syncthreadsor.cc + Atomic_func.cc ) if(UNIX) From 30d616d468e953588a66437d921572d1ad3cc0eb Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Tue, 7 Feb 2023 05:44:18 +0530 Subject: [PATCH 12/13] SWDEV-380433 - catch tests avoid same test detection (#157) fixes duplicate entry of tests in ctestfile Change-Id: Ia3e213109272b296e93e7b98656b080426e6ebb3 --- .../Catch2/cmake/Catch2/catch_include.cmake | 68 ++++++++++--------- 1 file changed, 35 insertions(+), 33 deletions(-) diff --git a/catch/external/Catch2/cmake/Catch2/catch_include.cmake b/catch/external/Catch2/cmake/Catch2/catch_include.cmake index b698e8b96c..9d37d62260 100644 --- a/catch/external/Catch2/cmake/Catch2/catch_include.cmake +++ b/catch/external/Catch2/cmake/Catch2/catch_include.cmake @@ -3,39 +3,41 @@ get_filename_component(_cmake_path cmake ABSOLUTE) -foreach(EXEC_NAME ${exc_names}) - if(WIN32) - set(EXEC_NAME ${EXEC_NAME}.exe) - endif() - if(EXISTS "${EXEC_NAME}") - execute_process( - COMMAND "${_cmake_path}" - -D "TEST_TARGET=${TARGET}" - -D "TEST_EXECUTABLE=${EXEC_NAME}" - -D "TEST_EXECUTOR=${crosscompiling_emulator}" - -D "TEST_WORKING_DIR=${_workdir}" - -D "TEST_SPEC=${_TEST_SPEC}" - -D "TEST_EXTRA_ARGS=${_EXTRA_ARGS}" - -D "TEST_PROPERTIES=${_PROPERTIES}" - -D "TEST_PREFIX=${_TEST_PREFIX}" - -D "TEST_SUFFIX=${_TEST_SUFFIX}" - -D "TEST_LIST=${_TEST_LIST}" - -D "TEST_REPORTER=${_REPORTER}" - -D "TEST_OUTPUT_DIR=${_OUTPUT_DIR}" - -D "TEST_OUTPUT_PREFIX=${_OUTPUT_PREFIX}" - -D "TEST_OUTPUT_SUFFIX=${_OUTPUT_SUFFIX}" - -D "CTEST_FILE=${ctestfilepath}" - -P "${_CATCH_ADD_TEST_SCRIPT}" - OUTPUT_VARIABLE output - RESULT_VARIABLE result - WORKING_DIRECTORY "${TEST_WORKING_DIR}" - ) - else() - message("executable not found : ${EXEC_NAME}" ) - endif() -endforeach() +if(NOT EXISTS "${ctestfilepath}") + foreach(EXEC_NAME ${exc_names}) + if(WIN32) + set(EXEC_NAME ${EXEC_NAME}.exe) + endif() + if(EXISTS "${EXEC_NAME}") + execute_process( + COMMAND "${_cmake_path}" + -D "TEST_TARGET=${TARGET}" + -D "TEST_EXECUTABLE=${EXEC_NAME}" + -D "TEST_EXECUTOR=${crosscompiling_emulator}" + -D "TEST_WORKING_DIR=${_workdir}" + -D "TEST_SPEC=${_TEST_SPEC}" + -D "TEST_EXTRA_ARGS=${_EXTRA_ARGS}" + -D "TEST_PROPERTIES=${_PROPERTIES}" + -D "TEST_PREFIX=${_TEST_PREFIX}" + -D "TEST_SUFFIX=${_TEST_SUFFIX}" + -D "TEST_LIST=${_TEST_LIST}" + -D "TEST_REPORTER=${_REPORTER}" + -D "TEST_OUTPUT_DIR=${_OUTPUT_DIR}" + -D "TEST_OUTPUT_PREFIX=${_OUTPUT_PREFIX}" + -D "TEST_OUTPUT_SUFFIX=${_OUTPUT_SUFFIX}" + -D "CTEST_FILE=${ctestfilepath}" + -P "${_CATCH_ADD_TEST_SCRIPT}" + OUTPUT_VARIABLE output + RESULT_VARIABLE result + WORKING_DIRECTORY "${TEST_WORKING_DIR}" + ) + else() + message("executable not found : ${EXEC_NAME}" ) + endif() + endforeach() +endif() if(EXISTS "${ctestfilepath}") -# include the generated ctest file for execution -include(${ctestfilepath}) + # include the generated ctest file for execution + include(${ctestfilepath}) endif() From 554ab5dc40af4ba73aabec6c372d4e3c52fd2cdb Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Tue, 7 Feb 2023 13:54:34 +0530 Subject: [PATCH 13/13] SWDEV-369555 - Update and enable hipMemAdvise test (#126) Change-Id: If291514eeacef21dea99bc4b4a78bc98fb0bc1ca --- catch/unit/memory/hipMemAdvise_old.cc | 15 ++++++--------- 1 file changed, 6 insertions(+), 9 deletions(-) diff --git a/catch/unit/memory/hipMemAdvise_old.cc b/catch/unit/memory/hipMemAdvise_old.cc index cef164e2ff..b951998053 100644 --- a/catch/unit/memory/hipMemAdvise_old.cc +++ b/catch/unit/memory/hipMemAdvise_old.cc @@ -232,15 +232,6 @@ TEST_CASE("Unit_hipMemAdvise_NegtveTsts") { std::string str; HIP_CHECK(hipGetDeviceCount(&NumDevs)); HIP_CHECK(hipMallocManaged(&Hmm, MEM_SIZE * 2, hipMemAttachGlobal)); -#if HT_AMD - // Passing invalid value(99) device param - IfTestPassed &= CheckError(hipMemAdvise(Hmm, MEM_SIZE * 2, - hipMemAdviseSetReadMostly, 99), __LINE__); - - // Passing invalid value(-12) device param - IfTestPassed &= CheckError(hipMemAdvise(Hmm, MEM_SIZE * 2, - hipMemAdviseSetReadMostly, -12), __LINE__); -#endif // Passing NULL as first parameter instead of valid pointer to a memory IfTestPassed &= CheckError(hipMemAdvise(NULL, MEM_SIZE * 2, hipMemAdviseSetReadMostly, 0), __LINE__); @@ -378,6 +369,12 @@ TEST_CASE("Unit_hipMemAdvise_ReadMostly") { WARN("out value: " << out); IfTestPassed = false; } + // hipMemAdvise should succeed for SetReadMostly and UnsetReadMostly + // irrespective of the device + HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseSetReadMostly, 99)); + + HIP_CHECK(hipMemAdvise(Hmm, MEM_SIZE, hipMemAdviseUnsetReadMostly, -12)); + HIP_CHECK(hipFree(Hmm)); REQUIRE(IfTestPassed); } else {