Enable check-openmp in AOMP.

The way to run the AOMP check-make tests is:

cd build/llvm-project
AOMP=~/rocm/aomp AOMP_GPU=gfx90a make check-openmp -j32

Change-Id: Icec3bb75cce3c166fc7b6c010028a27ca2fdd6ad

Author: doru1004

openmp/libomptarget/include/OpenMP/Mapping.h

@@ -72,8 +72,6 @@ struct HostDataToTargetTy {
   const uintptr_t TgtAllocBegin; // allocated target memory
   const uintptr_t TgtPtrBegin; // mapped target memory = TgtAllocBegin + padding
 
-  const bool IsUSMAlloc; // used to track maps under USM mode (optional)
-
 private:
   static const uint64_t INFRefCount = ~(uint64_t)0;
   static std::string refCountToStr(uint64_t RefCount) {
@@ -127,10 +125,9 @@ struct HostDataToTargetTy {
   HostDataToTargetTy(uintptr_t BP, uintptr_t B, uintptr_t E,
                      uintptr_t TgtAllocBegin, uintptr_t TgtPtrBegin,
                      bool UseHoldRefCount, map_var_info_t Name = nullptr,
-                     bool IsINF = false, bool IsUSMAlloc = false)
+                     bool IsINF = false)
       : HstPtrBase(BP), HstPtrBegin(B), HstPtrEnd(E), HstPtrName(Name),
         TgtAllocBegin(TgtAllocBegin), TgtPtrBegin(TgtPtrBegin),
-        IsUSMAlloc(IsUSMAlloc),
         States(std::make_unique<StatesTy>(UseHoldRefCount ? 0
                                           : IsINF         ? INFRefCount
                                                           : 1,

openmp/libomptarget/include/Shared/PluginAPI.h

@@ -42,8 +42,6 @@ bool __tgt_rtl_has_USM_capable_dGPU(void);
 
 bool __tgt_rtl_are_allocations_for_maps_on_apus_disabled(void);
 
-bool __tgt_rtl_is_no_maps_check(void);
-
 bool __tgt_rtl_is_fine_grained_memory_enabled(void);
 
 // Set up environement e.g. depending on the values of the env vars

openmp/libomptarget/include/Shared/PluginAPI.inc

@@ -53,7 +53,6 @@ PLUGIN_API_HANDLE(has_apu_device, true);
 PLUGIN_API_HANDLE(has_USM_capable_dGPU, true);
 PLUGIN_API_HANDLE(are_allocations_for_maps_on_apus_disabled, true);
 PLUGIN_API_HANDLE(requested_prepopulate_gpu_page_table, true);
-PLUGIN_API_HANDLE(is_no_maps_check, true);
 PLUGIN_API_HANDLE(is_fine_grained_memory_enabled, true);
 PLUGIN_API_HANDLE(is_system_supporting_managed_memory, true);
 PLUGIN_API_HANDLE(number_of_team_procs, true);

openmp/libomptarget/plugins-nextgen/amdgpu/src/rtl.cpp

@@ -848,9 +848,16 @@ struct AMDGPUKernelTy : public GenericKernelTy {
         ThreadLimitClause[0] += GenericDevice.getWarpSize();
     }
 
-    return std::min(MaxNumThreads, (ThreadLimitClause[0] > 0)
-                                       ? ThreadLimitClause[0]
-                                       : PreferredNumThreads);
+    // Limit number of threads taking into consideration the user
+    // environment variable OMP_TEAMS_THREAD_LIMIT if provided.
+    uint32_t CurrentMaxNumThreads = MaxNumThreads;
+    if (TeamsThreadLimitEnvVar > 0)
+      CurrentMaxNumThreads = std::min(
+          static_cast<uint32_t>(TeamsThreadLimitEnvVar), CurrentMaxNumThreads);
+
+    return std::min(CurrentMaxNumThreads,
+                    (ThreadLimitClause[0] > 0) ? ThreadLimitClause[0] :
+                    PreferredNumThreads);
   }
   uint64_t getNumBlocks(GenericDeviceTy &GenericDevice,
                         uint32_t NumTeamsClause[3], uint64_t LoopTripCount,
@@ -4026,8 +4033,6 @@ struct AMDGPUPluginTy final : public GenericPluginTy {
     return PrepopulateGPUPageTable;
   }
 
-  bool IsNoMapsCheck() override final { return NoUSMMapChecks; }
-
   bool IsFineGrainedMemoryEnabled() override final {
     return EnableFineGrainedMemory;
   }
@@ -4047,19 +4052,13 @@ struct AMDGPUPluginTy final : public GenericPluginTy {
     if (!Initialized)
       FATAL_MESSAGE(1, "%s", "parseEnvVars was called on uninitialized plugin");
 
-    NoMapChecks = BoolEnvar("OMPX_DISABLE_MAPS", true);
     DisableUsmMaps = BoolEnvar("OMPX_DISABLE_USM_MAPS", false);
     HsaXnack = BoolEnvar("HSA_XNACK", false);
     APUPrefault = BoolEnvar("OMPX_EAGER_ZERO_COPY_MAPS", false);
     ZeroCopyForMapsOnUsm = BoolEnvar("OMPX_APU_MAPS", false);
   }
 
   void setUpEnv() override final {
-
-    if (NoMapChecks.get() == false) {
-      NoUSMMapChecks = false;
-    }
-
     if (DisableUsmMaps.get() == true) {
       EnableFineGrainedMemory = true;
     }
@@ -4360,10 +4359,6 @@ struct AMDGPUPluginTy final : public GenericPluginTy {
   // page table.
   bool PrepopulateGPUPageTable{false};
 
-  // Set by OMPX_DISABLE_MAPS environment variable.
-  // When active (default value), maps are ignored by the runtime
-  bool NoUSMMapChecks{true};
-
   // Set by OMPX_DISABLE_USM_MAPS environment variable.
   // If set, fine graned memory is used for maps instead of coarse grained.
   bool EnableFineGrainedMemory{false};

openmp/libomptarget/plugins-nextgen/common/include/PluginInterface.h

@@ -1104,8 +1104,6 @@ struct GenericPluginTy {
 
   virtual bool requestedPrepopulateGPUPageTable() { return false; }
 
-  virtual bool IsNoMapsCheck() { return false; }
-
   virtual bool IsFineGrainedMemoryEnabled() { return false; }
 
   virtual bool IsSystemSupportingManagedMemory() { return false; }

openmp/libomptarget/plugins-nextgen/common/src/PluginInterface.cpp

@@ -1853,8 +1853,6 @@ bool __tgt_rtl_requested_prepopulate_gpu_page_table() {
   return Plugin::get().requestedPrepopulateGPUPageTable();
 }
 
-bool __tgt_rtl_is_no_maps_check() { return Plugin::get().IsNoMapsCheck(); }
-
 bool __tgt_rtl_is_fine_grained_memory_enabled() {
   return Plugin::get().IsFineGrainedMemoryEnabled();
 }
@@ -1935,7 +1933,8 @@ void *__tgt_rtl_data_alloc(int32_t DeviceId, int64_t Size, void *HostPtr,
   assert(*AllocOrErr && "Null pointer upon successful allocation");
 
   // Method has no effect when the CUDA Plugin is used.
-  if (Kind == TARGET_ALLOC_SHARED)
+  // This method can only be called if HostPtr is not null.
+  if (HostPtr && Kind == TARGET_ALLOC_SHARED)
     __tgt_rtl_set_coarse_grain_mem_region(DeviceId, HostPtr, Size);
 
   return *AllocOrErr;

openmp/libomptarget/plugins-nextgen/common/src/trace.h

@@ -350,16 +350,6 @@ bool __tgt_rtl_is_fine_grained_memory_enabled() {
 #define __tgt_rtl_is_fine_grained_memory_enabled(...)                          \
   __tgt_rtl_is_fine_grained_memory_enabled_impl(__VA_ARGS__)
 
-static bool __tgt_rtl_is_no_maps_check_impl();
-bool __tgt_rtl_is_no_maps_check() {
-  auto t = detail::log<bool>(__func__);
-  bool r = __tgt_rtl_is_no_maps_check_impl();
-  t.res(r);
-  return r;
-}
-#define __tgt_rtl_is_no_maps_check(...)                                        \
-  __tgt_rtl_is_no_maps_check_impl(__VA_ARGS__)
-
 static int32_t __tgt_rtl_launch_kernel_sync_impl(int32_t device_id,
                                                  void *tgt_entry_ptr,
                                                  void **tgt_args,

openmp/libomptarget/src/OpenMP/Mapping.cpp

@@ -280,19 +280,6 @@ TargetPointerResultTy MappingInfoTy::getTargetPointer(
           Device.RTL->prepopulate_page_table) {
         Device.RTL->prepopulate_page_table(Device.DeviceID, HstPtrBegin, Size);
       }
-
-      if (!Device.RTL->is_no_maps_check()) {
-        // even under unified_shared_memory need to check for correctness of
-        // use of map clauses. Device pointer is same as host ptr in this case
-        LR.TPR.setEntry(HDTTMap
-                            ->emplace(new HostDataToTargetTy(
-                                (uintptr_t)HstPtrBase, (uintptr_t)HstPtrBegin,
-                                (uintptr_t)HstPtrBegin + Size,
-                                (uintptr_t)HstPtrBegin, (uintptr_t)HstPtrBegin,
-                                HasHoldModifier, HstPtrName, /*IsInf=*/true,
-                                /*IsUSMAlloc=*/true))
-                            .first->HDTT);
-      }
     }
     DP("Return HstPtrBegin " DPxMOD " Size=%" PRId64 " for unified shared "
        "memory\n",
@@ -402,8 +389,7 @@ TargetPointerResultTy MappingInfoTy::getTgtPtrBegin(
   LR.TPR.Flags.IsPresent = true;
 
   if ((LR.Flags.IsContained ||
-       (!MustContain && (LR.Flags.ExtendsBefore || LR.Flags.ExtendsAfter))) &&
-      !LR.TPR.getEntry()->IsUSMAlloc) {
+       (!MustContain && (LR.Flags.ExtendsBefore || LR.Flags.ExtendsAfter)))) {
 
     LR.TPR.Flags.IsLast =
         LR.TPR.getEntry()->decShouldRemove(UseHoldRefCount, ForceDelete);

openmp/libomptarget/test/api/amd_assert.c

@@ -0,0 +1,15 @@
+// RUN: %libomptarget-compile-generic -O0 && %libomptarget-run-generic 2>&1 | %fcheck-generic
+// RUN: %libomptarget-compileopt-run-and-check-generic
+
+#include <assert.h>
+#include <stdio.h>
+
+int main() {
+  int i = 1;
+#pragma omp target
+  assert(i > 0);
+
+  // CHECK: PASS
+  printf("PASS\n");
+  return 0;
+}

openmp/libomptarget/test/api/assert.c

@@ -1,3 +1,4 @@
+// XFAIL: amdgcn-amd-amdhsa
 // RUN: %libomptarget-compile-run-and-check-generic
 // RUN: %libomptarget-compileopt-run-and-check-generic