[AMDGPU][Offload] Adjustment to reflect upstream PR151882

Author: hidekisaito

offload/plugins-nextgen/amdgpu/src/rtl.cpp

@@ -4513,36 +4513,34 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
 
   bool checkIfCoarseGrainMemoryNearOrAbove64GB() {
     for (AMDGPUMemoryPoolTy *Pool : AllMemoryPools) {
-      if (Pool->isGlobal() && Pool->isCoarseGrained()) {
-	uint64_t Value;
-        hsa_status_t Status =
-            Pool->getAttrRaw(HSA_AMD_MEMORY_POOL_INFO_SIZE, Value);
-        if (Status != HSA_STATUS_SUCCESS) continue;
-	constexpr uint64_t Almost64Gig = 0xFF0000000;
-	if (Value >= Almost64Gig) return true;
-      }
+      if (!Pool->isGlobal() || !Pool->isCoarseGrained())
+        continue;
+      uint64_t Value;
+      hsa_status_t Status =
+          Pool->getAttrRaw(HSA_AMD_MEMORY_POOL_INFO_SIZE, Value);
+      if (Status != HSA_STATUS_SUCCESS)
+        continue;
+      constexpr uint64_t Almost64Gig = 0xFF0000000;
+      if (Value >= Almost64Gig)
+        return true;
     }
     return false; // CoarseGrain pool w/ 64GB or more capacity not found
   }
 
   size_t getMemoryManagerSizeThreshold() override {
-    // TODO: check performance on lower memory capacity GPU
-    // for lowering the threshold from 64GB.
+    // Targeting high memory capacity GPUs such as
+    // data center GPUs.
     if (checkIfCoarseGrainMemoryNearOrAbove64GB()) {
-      // Set GenericDeviceTy::MemoryManager's Threshold to ~2GB,
-      // used if not set by LIBOMPTARGET_MEMORY_MANAGER_THRESHOLD
-      // ENV var. This MemoryManager is used for
-      // omp_target_alloc(), OpenMP (non-usm) map clause, etc.
-      //
-      // TODO 1: Fine tune to lower the threshold closer to 1GB.
-      // TODO 2: HSA-level memory manager on the user-side such that
-      // memory management is shared with HIP and OpenCL.
+      // Set GenericDeviceTy::MemoryManager's Threshold to 3GiB,
+      // if threshold is not already set by ENV var
+      // LIBOMPTARGET_MEMORY_MANAGER_THRESHOLD.
+      // This MemoryManager is used for omp_target_alloc(), OpenMP
+      // (non-usm) map clause, etc.
       //
-      // If this value needs to go above UINT_MAX, consider
-      // adding sizeof(size_t) check to avoid unpleasant truncation
-      // surprises where size_t is still 32bit.
-      constexpr size_t Almost3Gig = 3000000000u;
-      return Almost3Gig;
+      // Ideally, this kind of pooling is best performed at
+      // a common level (e.g, user side of HSA) between OpenMP and HIP
+      // but that feature does not exist (yet).
+      return 3ul * 1024 * 1024 * 1024 /* 3 GiB */;
     }
     return 0;
   }

offload/plugins-nextgen/common/include/PluginInterface.h

@@ -1342,7 +1342,7 @@ struct GenericDeviceTy : public DeviceAllocatorTy {
   MemoryManagerTy *MemoryManager;
 
   /// Per device setting of MemoryManager's Threshold
-  virtual size_t getMemoryManagerSizeThreshold() { return 0 /* use default */; }
+  virtual size_t getMemoryManagerSizeThreshold() { return 0; }
 
   /// Environment variables defined by the OpenMP standard.
   Int32Envar OMP_TeamLimit;

offload/test/lit.cfg

@@ -121,6 +121,7 @@ if config.libomptarget_test_pgo:
 # For all other targets, we currently assume it is.
 supports_unified_shared_memory = True
 supports_apu = False
+supports_large_allocation_memory_pool = False
 is_mi200 = False
 if config.libomptarget_current_target.startswith('nvptx'):
   try:
@@ -133,9 +134,11 @@ if config.libomptarget_current_target.startswith('nvptx'):
 elif config.libomptarget_current_target.startswith('amdgcn'):
     # amdgpu_test_arch contains a list of AMD GPUs in the system
     # only check the first one assuming that we will run the test on it.
-    if not (config.amdgpu_test_arch.startswith("gfx90a") or
-            config.amdgpu_test_arch.startswith("gfx942") or
-            config.amdgpu_test_arch.startswith("gfx950")):
+    if (config.amdgpu_test_arch.startswith("gfx90a") or
+        config.amdgpu_test_arch.startswith("gfx942") or
+        config.amdgpu_test_arch.startswith("gfx950")):
+       supports_large_allocation_memory_pool = True
+    else:
        supports_unified_shared_memory = False
     # check if AMD architecture is an APU:
     if ((config.amdgpu_test_arch.startswith("gfx942") and
@@ -152,7 +155,8 @@ elif config.libomptarget_current_target.startswith('amdgcn'):
         config.available_features.add('apu')
     if is_mi200:
         config.available_features.add('mi200')
-
+if supports_large_allocation_memory_pool:
+   config.available_features.add('large_allocation_memory_pool')
 
 # Setup environment to find dynamic library at runtime
 if config.operating_system == 'Windows':

offload/test/sanitizer/use_after_free_2.c

@@ -10,6 +10,10 @@
 // UNSUPPORTED: s390x-ibm-linux-gnu
 // UNSUPPORTED: s390x-ibm-linux-gnu-LTO
 
+// If offload memory pooling is enabled for a large allocation, reuse error is
+// not detected.
+// UNSUPPORTED: large_allocation_memory_pool
+
 #include <omp.h>
 
 int main() {

offload/test/sanitizer/use_after_free_3.c

@@ -0,0 +1,38 @@
+// clang-format off
+// RUN: %libomptarget-compileopt-generic
+// RUN: %not --crash env -u LLVM_DISABLE_SYMBOLIZATION OFFLOAD_TRACK_ALLOCATION_TRACES=1 LIBOMPTARGET_MEMORY_MANAGER_THRESHOLD=1024 %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=CHECK
+// RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic --check-prefixes=CHECK-PASS
+// clang-format on
+
+// If offload memory pooling is enabled for a large allocation, reuse error is
+// not detected. Run the test w/ and w/o ENV var override on memory pooling
+// threshold.
+// REQUIRES: large_allocation_memory_pool
+
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+  int N = (1 << 30);
+  char *A = (char *)malloc(N);
+  char *P;
+#pragma omp target map(A[ : N]) map(from : P)
+  {
+    P = &A[N / 2];
+    *P = 3;
+  }
+  // clang-format off
+// CHECK: OFFLOAD ERROR: memory access fault by GPU {{.*}} (agent 0x{{.*}}) at virtual address [[PTR:0x[0-9a-z]*]]. Reasons: {{.*}}
+// CHECK: Device pointer [[PTR]] points into prior host-issued allocation:
+// CHECK: Last deallocation:
+// CHECK: Last allocation of size 1073741824
+// clang-format on
+#pragma omp target
+  {
+    *P = 5;
+  }
+
+  // CHECK-PASS: PASS
+  printf("PASS\n");
+  return 0;
+}