SWDEV-413161 - [catch2][dtest] hip directed catch2 with xnack+ test added for hipMemAdvise api for TstAlignedAllocMem

Change-Id: I51252dfff34f15467d56cefc920feb300868df13


[ROCm/hip-tests commit: 658123244a]

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

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 <stdlib.h>
+#include <iostream>
+#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<int*>(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

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 <hip_test_common.hh>
 #include <hip_test_features.hh>
+#include <hip_test_process.hh>
+
 #if __linux__
 #include <unistd.h>
 #include <sys/mman.h>
@@ -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) {
+      for (int i = 0; i < NumElms; ++i) {
-          if (Mllc[i] != (InitVal + 10)) {
+        if (Mllc[i] != (InitVal + 10)) {
-            DataMismatch++;
+          DataMismatch++;
-          }
         }
-      REQUIRE(DataMismatch == 0);
       }
-  } else {
+      REQUIRE(DataMismatch == 0);
-      HipTest::HIP_SKIP_TEST("GPU is not xnack enabled hence skipping the test...\n");
     }
+  } 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");
   }
 }
-