[OpenMP] Enable map checks under unified_shared_memory mode

[carlo.bertolli]

Summary

Enable correctness checks in runtime for map clauses under unified_shared_memory mode.
Add OPENMP_DISABLE_MAPS environment variable input to disable this behavior.

As an example, before this patch is applied, this example would not fail:

#pragma omp requires unified_shared_memory

void foo() {

int a[10];

#pragma omp target enter data map(to:a[0:5])
#pragma omp target enter data map(to:a[0:10]) // extending already mapped memory, illegal!

}

This is illegal in OpenMP, and presumably also under unified_shared_memory mode, but the current implementation bails out on map checks under this mode.
By specifications, under unified_shared_memory mode, the user is allowed not to use maps and can expect the program to work on an system that supports the mode. However, for portability to systems that do not support unified_shared_memory, map clauses should always be correct. Not applying checks in the runtime should be a conscious request by the user.

[carlo.bertolli]

Update environment variable name to reflect comments and intended name

[jdoerfert]

As discussed, we should put this under an environment variable switch and do some measurements.

openmp/libomptarget/src/rtl.h

119

This is potentially uninitialized.

[grokos]

openmp/libomptarget/src/device.cpp

236–238

We will have a problem here after D107925, D107927 and D107928 land. These patches assume that USM mappings do not have an entry in HostDataToTargetMap. That's how they figure out whether we have an allocation in shared memory or in GPU memory, i.e. the value of IsHostPtr is determined by the presence or absence of an entry in the map. If we want to move on with this change here, we'll need to introduce an extra member in struct HostDataToTargetTy, e.g. bool IsUSMAlloc = false;, and this emplacement here will have to be extended with ..., HstPtrName, /*IsUSMAlloc*/ true). In general, I think that implementing desired support involves much more work than this patch. Basically, we need to treat USM data the exact same way as device-side mappings, i.e. remove all specialized code paths for USM allocations and only check whether we have USM data when we are about to alloc memory / perform data transfers.

370–372

Rebase this patch after D107925 lands, as these lines will be removed.

387–390

This check can go away after D107925 lands because the caller of this function has already performed this check.

openmp/libomptarget/src/rtl.cpp

82

It's better to rename this env var to something like LIBOMPTARGET_DISABLE_USM_MAPS. All env vars begin either with OMP_ if they are part of the OpenMP standard or with LIBOMPTARGET_ if they are libomptarget-specific extensions. We have no env vars starting with OPENMP_, so let's keep it consistent with the existing codebase. Also, including "USM" in the name makes it more descriptive of exactly what it does.

openmp/libomptarget/src/rtl.h

119

Need to initialize, I guess we want the default value to be false?

Also, can you add a comment about what this variable does? And maybe a more meaningful name would be something like "NoUSMMapEntries"?

[carlo.bertolli]

Updated patch based on comments: add new map table entry field to track USM maps; use it to determine behavior in getTargetPointer and deallocTgtPtr; update var name and add explaining comments.

openmp/libomptarget/src/device.cpp

if (LR.Flags.IsContained ||
     // host address where correctness is not impeded. In all other cases maps
     // are respected.
     // In addition to the mapping rules above, the close map modifier forces the
     // mapping of the variable to the device.
     if (Size) {
       DP("Return HstPtrBegin " DPxMOD " Size=%" PRId64 " for unified shared "
          "memory\n",
          DPxPTR((uintptr_t)HstPtrBegin), Size);
+
+      if (!PM->RTLs.NoMaps)
+        // 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
+        HostDataToTargetMap.emplace(
+            (uintptr_t)HstPtrBase, (uintptr_t)HstPtrBegin,
+            (uintptr_t)HstPtrBegin + Size, (uintptr_t)HstPtrBegin, HstPtrName);
+
       IsHostPtr = true;
       TargetPointer = HstPtrBegin;
     }
   } else if (HasPresentModifier) {
     DP("Mapping required by 'present' map type modifier does not exist for "
        "HstPtrBegin=" DPxMOD ", Size=%" PRId64 "\n",
        DPxPTR(HstPtrBegin), Size);
     MESSAGE("device mapping required by 'present' map type modifier does not "
if (lr.Flags.IsContained || lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter) {
     return (void *)tp;
   }
 
   return NULL;
 }
 
 int DeviceTy::deallocTgtPtr(void *HstPtrBegin, int64_t Size,
                             bool HasCloseModifier) {
-  if (PM->RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY &&
-      !HasCloseModifier)
-    return OFFLOAD_SUCCESS;
   // Check if the pointer is contained in any sub-nodes.
   int rc;
   DataMapMtx.lock();
   LookupResult lr = lookupMapping(HstPtrBegin, Size);
   if (lr.Flags.IsContained || lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter) {
     auto &HT = *lr.Entry;
     if (HT.decRefCount() == 0) {
       DP("Deleting tgt data " DPxMOD " of size %" PRId64 "\n",
          DPxPTR(HT.TgtPtrBegin), Size);
+      if (!(PM->RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY &&
+            !HasCloseModifier))
         deleteData((void *)HT.TgtPtrBegin);
       INFO(OMP_INFOTYPE_MAPPING_CHANGED, DeviceID,
            "Removing map entry with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD
            ", Size=%" PRId64 ", Name=%s\n",
            DPxPTR(HT.HstPtrBegin), DPxPTR(HT.TgtPtrBegin), Size,
            (HT.HstPtrName) ? getNameFromMapping(HT.HstPtrName).c_str()
                            : "unknown");
+      if (!PM->RTLs.NoMaps)
         HostDataToTargetMap.erase(lr.Entry);
     }
     rc = OFFLOAD_SUCCESS;
   } else {
     REPORT("Section to delete (hst addr " DPxMOD ") does not exist in the"
            " allocated memory\n",
            DPxPTR(HstPtrBegin));
     rc = OFFLOAD_FAIL;
   }

openmp/libomptarget/src/rtl.h

struct RTLsTy {
   std::list<RTLInfoTy> AllRTLs;
 
   // Array of pointers to the detected runtime libraries that have compatible
   // binaries.
   std::vector<RTLInfoTy *> UsedRTLs;
 
   int64_t RequiresFlags = OMP_REQ_UNDEFINED;
 
+  bool NoMaps;
+
   explicit RTLsTy() = default;
 
   // Register the clauses of the requires directive.
   void RegisterRequires(int64_t flags);
 
   // Initialize RTL if it has not been initialized
   void initRTLonce(RTLInfoTy &RTL);
 

openmp/libomptarget/src/rtl.cpp

 void RTLsTy::LoadRTLs() {
   // Parse environment variable OMP_TARGET_OFFLOAD (if set)
   PM->TargetOffloadPolicy =
       (kmp_target_offload_kind_t)__kmpc_get_target_offload();
   if (PM->TargetOffloadPolicy == tgt_disabled) {
     return;
   }
 
+  // In UNIFIED_SHARED_MEMORY mode, users can disable
+  // map checks by setting OPENMP_NO_MAPS
+  if (const char *NoMapsStr = getenv("OPENMP_NO_MAPS"))
+    if (NoMapsStr)
+      NoMaps = std::stoi(NoMapsStr);
+
   DP("Loading RTLs...\n");
 
   // Attempt to open all the plugins and, if they exist, check if the interface
   // is correct and if they are supporting any devices.
   for (auto *Name : RTLNames) {
     DP("Loading library '%s'...\n", Name);
     void *dynlib_handle = dlopen(Name, RTLD_NOW);
 

openmp/libomptarget/test/unified_shared_memory/check_maps.c

+// RUN: %libomptarget-compile-generic
+// RUN: %libomptarget-run-fail-generic 2>&1 \
+// RUN: | %fcheck-generic
+
+// CHECK: Libomptarget message: explicit extension not allowed: host address specified is 0x{{.*}}, but device allocation maps to host at 0x{{.*}}
+// CHECK: Libomptarget error: Call to getOrAllocTgtPtr returned null pointer (device failure or illegal mapping).
+// CHECK: Libomptarget fatal error 1: failure of target construct while offloading is mandatory
+
+// REQUIRES: unified_shared_memory
+
+#include <stdio.h>
+#include <omp.h>
+
+// ---------------------------------------------------------------------------
+// Various definitions copied from OpenMP RTL
+
+extern void __tgt_register_requires(int64_t);
+
+// End of definitions copied from OpenMP RTL.
+// ---------------------------------------------------------------------------
+
+#pragma omp requires unified_shared_memory
+
+#define N 1024
+
+int main(int argc, char *argv[]) {
+  int data[N];
+
+  // Manual registration of requires flags for Clang versions
+  // that do not support requires.
+  __tgt_register_requires(8);
+
+
+#pragma omp target enter data map(to:data[0:N/2])
+
+// extending a map under unified_shared_memory mode is prohibited
+#pragma omp target enter data map(to:data[0:N])
+
+#pragma omp target
+  {
+  }
+
+  return 0;
+}