From d5f76714e87dfd7190ecb906bc9bbf31f0c1cf6b Mon Sep 17 00:00:00 2001 From: mbhiutra Date: Fri, 29 Dec 2023 19:39:49 +0530 Subject: [PATCH] SWDEV-413161 - [catch2][dtest] hip directed catch2 with xnack+ test added for hipMemAdvise api for TstAlignedAllocMem Change-Id: I51252dfff34f15467d56cefc920feb300868df13 [ROCm/hip-tests commit: 658123244ae7ae7902167f5dda32e93c7a19f3de] --- .../catch/unit/memory/CMakeLists.txt | 4 + .../hipMemAdvise_AlignedAllocMem_Exe.cc | 85 +++++++++++++++++++ .../catch/unit/memory/hipMemAdvise_old.cc | 48 ++++++++--- 3 files changed, 124 insertions(+), 13 deletions(-) create mode 100644 projects/hip-tests/catch/unit/memory/hipMemAdvise_AlignedAllocMem_Exe.cc diff --git a/projects/hip-tests/catch/unit/memory/CMakeLists.txt b/projects/hip-tests/catch/unit/memory/CMakeLists.txt index 025e369b50..5dd1554e7f 100644 --- a/projects/hip-tests/catch/unit/memory/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/memory/CMakeLists.txt @@ -111,6 +111,10 @@ hip_add_exe_to_target(NAME MemoryTest1 if(HIP_PLATFORM MATCHES "amd") set_source_files_properties(hipHostRegister.cc PROPERTIES COMPILE_FLAGS -std=c++17) add_executable(hipHostRegisterPerf EXCLUDE_FROM_ALL hipHostRegister_exe.cc) +if(UNIX) + add_executable(hipMemAdviseTstAlignedAllocMem EXCLUDE_FROM_ALL hipMemAdvise_AlignedAllocMem_Exe.cc) + add_dependencies(MemoryTest1 hipMemAdviseTstAlignedAllocMem) +endif() endif() set(TEST_SRC diff --git a/projects/hip-tests/catch/unit/memory/hipMemAdvise_AlignedAllocMem_Exe.cc b/projects/hip-tests/catch/unit/memory/hipMemAdvise_AlignedAllocMem_Exe.cc new file mode 100644 index 0000000000..0fe16a273b --- /dev/null +++ b/projects/hip-tests/catch/unit/memory/hipMemAdvise_AlignedAllocMem_Exe.cc @@ -0,0 +1,85 @@ +/* Copyright (c) 2024 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. +*/ + +#if __linux__ +#include +#include +#include "hip/hip_runtime_api.h" + +#define HIP_CHECK(error) \ +{ \ + hipError_t localError = error; \ + if ((localError != hipSuccess) && \ + (localError != hipErrorPeerAccessAlreadyEnabled)) { \ + return -1; \ + } \ +} + +// Kernel +__global__ void MemAdvise_Exe(int *Hmm, int n) { + for (int i = 0; i < n; i ++) { + Hmm[i] = Hmm[i] + 10; + } +} + +static int hipMemAdvise_AlignedAllocMem_Exe() { + int managedMem = 0, pageMemAccess = 0; + HIP_CHECK(hipDeviceGetAttribute(&pageMemAccess, + hipDeviceAttributePageableMemoryAccess, 0)); + std::cout << "\n hipDeviceAttributePageableMemoryAccess:" << pageMemAccess; + HIP_CHECK(hipDeviceGetAttribute(&managedMem, + hipDeviceAttributeManagedMemory, 0)); + std::cout << "\n 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)); + hipError_t err = hipMemPrefetchAsync(Mllc, MemSz, 0, strm); + if (err != hipErrorInvalidValue) return -1; + HIP_CHECK(hipStreamSynchronize(strm)); + MemAdvise_Exe<<<(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) return -1; + } + return 0; +} + +int main() { + return hipMemAdvise_AlignedAllocMem_Exe(); +} +#endif diff --git a/projects/hip-tests/catch/unit/memory/hipMemAdvise_old.cc b/projects/hip-tests/catch/unit/memory/hipMemAdvise_old.cc index 625c216ef6..fcf98258fb 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemAdvise_old.cc +++ b/projects/hip-tests/catch/unit/memory/hipMemAdvise_old.cc @@ -1,5 +1,5 @@ /* -Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2021-2024 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 @@ -68,6 +68,8 @@ THE SOFTWARE. #include #include +#include + #if __linux__ #include #include @@ -373,9 +375,8 @@ TEST_CASE("Unit_hipMemAdvise_ReadMostly") { // 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 { @@ -670,7 +671,8 @@ TEST_CASE("Unit_hipMemAdvise_TstAlignedAllocMem") { HIP_CHECK(hipDeviceGetAttribute(&pageMemAccess, hipDeviceAttributePageableMemoryAccess, 0)); WARN("hipDeviceAttributePageableMemoryAccess:" << pageMemAccess); - HIP_CHECK(hipDeviceGetAttribute(&managedMem, hipDeviceAttributeManagedMemory, 0)); + HIP_CHECK(hipDeviceGetAttribute(&managedMem, + hipDeviceAttributeManagedMemory, 0)); WARN("hipDeviceAttributeManagedMemory: " << managedMem); if ((managedMem == 1) && (pageMemAccess == 1)) { int *Mllc = nullptr, MemSz = 4096 * 4, NumElms = 4096, InitVal = 123; @@ -685,20 +687,41 @@ TEST_CASE("Unit_hipMemAdvise_TstAlignedAllocMem") { // 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_ERROR(hipMemPrefetchAsync(Mllc, MemSz, 0, strm), hipErrorInvalidValue); + HIP_CHECK_ERROR(hipMemPrefetchAsync(Mllc, MemSz, 0, strm), + hipErrorInvalidValue); 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++; - } + for (int i = 0; i < NumElms; ++i) { + if (Mllc[i] != (InitVal + 10)) { + DataMismatch++; } - REQUIRE(DataMismatch == 0); } - } else { - HipTest::HIP_SKIP_TEST("GPU is not xnack enabled hence skipping the test...\n"); + REQUIRE(DataMismatch == 0); } + } else { + HipTest::HIP_SKIP_TEST("GPU is not xnack enabled hence skipping the test"); + } +} + +TEST_CASE("Unit_hipMemAdvise_TstAlignedAllocMem_XNACK") { + if (setenv("HSA_XNACK", "1", 1) != 0) { + HipTest::HIP_SKIP_TEST("Unable to set xnack on environment variable."); + return; + } + + hipDeviceProp_t prop; + int device; + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipGetDeviceProperties(&prop, device)); + std::string gfxName(prop.gcnArchName); + + if (gfxName.find("xnack+") != std::string::npos) { + hip::SpawnProc proc("hipMemAdviseTstAlignedAllocMem", true); + REQUIRE(proc.run() == 0); + } else { + HipTest::HIP_SKIP_TEST("GPU is not xnack enabled hence skipping the test"); + } } #endif @@ -902,4 +925,3 @@ TEST_CASE("Unit_hipMemAdvise_TstSetUnsetPrfrdLoc") { "attribute. Hence skipping the testing with Pass result.\n"); } } -