[OpenACC] Prototype OpenMP extension: hold map type modifier

The `hold` map type modifier is an original OpenMP extension for the
sake of supporting OpenACC's two reference counters for tracking
device allocations.  In future commits, `hold` will be used where
OpenACC's structured reference counter is required (for data and
compute constructs), and it will be omitted where OpenACC's dynamic
reference counter is required (for enter data and exit data
directives).

Currently, `hold` is not thoroughly tested outside of translations
from OpenACC to OpenMP, so it is not yet recommended for general use
in OpenMP code.  This commit introduces some testing at the OpenMP
level, but missing testing includes at least Clang codegen and
`PTR_AND_OBJ` map entries (see todos in commit).

Author: jdenny-ornl

clang/include/clang/AST/OpenMPClause.h

@@ -5395,7 +5395,7 @@ class OMPMapClause final : public OMPMappableExprListClause<OMPMapClause>,
   OpenMPMapModifierKind MapTypeModifiers[NumberOfOMPMapClauseModifiers] = {
       OMPC_MAP_MODIFIER_unknown, OMPC_MAP_MODIFIER_unknown,
       OMPC_MAP_MODIFIER_unknown, OMPC_MAP_MODIFIER_unknown,
-      OMPC_MAP_MODIFIER_unknown};
+      OMPC_MAP_MODIFIER_unknown, OMPC_MAP_MODIFIER_unknown};
 
   /// Location of map-type-modifiers for the 'map' clause.
   SourceLocation MapTypeModifiersLoc[NumberOfOMPMapClauseModifiers];

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -10139,6 +10139,8 @@ def err_omp_map_shared_storage : Error<
   "variable already marked as mapped in current construct">;
 def err_omp_invalid_map_type_for_directive : Error<
   "%select{map type '%1' is not allowed|map type must be specified}0 for '#pragma omp %2'">;
+def err_omp_invalid_map_type_modifier_for_directive : Error<
+  "map type modifier '%0' is not allowed for '#pragma omp %1'">;
 def err_omp_no_clause_for_directive : Error<
   "expected at least one %0 clause for '#pragma omp %1'">;
 def err_omp_threadprivate_in_clause : Error<

clang/include/clang/Basic/OpenMPKinds.def

@@ -126,6 +126,10 @@ OPENMP_MAP_MODIFIER_KIND(present)
 // TODO: Currently, it is not well tested outside of translations from OpenACC
 // to OpenMP, so it is not yet recommended for general use in OpenMP code.
 OPENMP_MAP_MODIFIER_KIND(no_alloc)
+// This is an OpenMP extension for the sake of OpenACC support.
+// TODO: Currently, it is not well tested outside of translations from OpenACC
+// to OpenMP, so it is not yet recommended for general use in OpenMP code.
+OPENMP_MAP_MODIFIER_KIND(hold)
 
 // Modifiers for 'to' or 'from' clause.
 OPENMP_MOTION_MODIFIER_KIND(mapper)

clang/lib/CodeGen/CGOpenMPRuntime.cpp

@@ -7012,6 +7012,16 @@ class MappableExprsHandler {
     // OpenACC to OpenMP, so it is not yet recommended for general use in
     // OpenMP code.
     OMP_MAP_NO_ALLOC = 0x2000,
+    // Increment and decrement a separate reference counter so that the data
+    // cannot be deallocated within the associated region.  Thus, this flag is
+    // intended to be used on target and target data directives.  It is not
+    // intended to be used on target enter/exit data directives because they are
+    // inherently dynamic not structured.
+    // This is an OpenMP extension for the sake of OpenACC support.
+    // TODO: Currently, it is not well tested outside of translations from
+    // OpenACC to OpenMP, so it is not yet recommended for general use in
+    // OpenMP code.
+    OMP_MAP_HOLD = 0x4000,
     /// The 16 MSBs of the flags indicate whether the entry is member of some
     /// struct/class.
     OMP_MAP_MEMBER_OF = 0xffff000000000000,
@@ -7295,6 +7305,8 @@ class MappableExprsHandler {
     if (llvm::find(MapModifiers, OMPC_MAP_MODIFIER_no_alloc)
         != MapModifiers.end())
       Bits |= OMP_MAP_NO_ALLOC;
+    if (llvm::find(MapModifiers, OMPC_MAP_MODIFIER_hold) != MapModifiers.end())
+      Bits |= OMP_MAP_HOLD;
     if (llvm::find(MotionModifiers, OMPC_MOTION_MODIFIER_present)
         != MotionModifiers.end())
       Bits |= OMP_MAP_PRESENT;

clang/lib/Parse/ParseOpenMP.cpp

@@ -3127,7 +3127,8 @@ bool Parser::parseMapTypeModifiers(OpenMPVarListDataTy &Data) {
     if (TypeModifier == OMPC_MAP_MODIFIER_always ||
         TypeModifier == OMPC_MAP_MODIFIER_close ||
         TypeModifier == OMPC_MAP_MODIFIER_present ||
-        TypeModifier == OMPC_MAP_MODIFIER_no_alloc) {
+        TypeModifier == OMPC_MAP_MODIFIER_no_alloc ||
+        TypeModifier == OMPC_MAP_MODIFIER_hold) {
       Data.MapTypeModifiers.push_back(TypeModifier);
       Data.MapTypeModifiersLoc.push_back(Tok.getLocation());
       ConsumeToken();

clang/lib/Sema/SemaOpenMP.cpp

@@ -17460,6 +17460,10 @@ static void checkMappableExpressionList(
   bool UpdateUMIt = false;
   Expr *UnresolvedMapper = nullptr;
 
+  bool HasHoldModifier =
+      Modifiers.end() !=
+      std::find(Modifiers.begin(), Modifiers.end(), OMPC_MAP_MODIFIER_hold);
+
   // Keep track of the mappable components and base declarations in this clause.
   // Each entry in the list is going to have a list of components associated. We
   // record each set of the components so that we can build the clause later on.
@@ -17659,6 +17663,19 @@ static void checkMappableExpressionList(
         continue;
       }
 
+      // The 'hold' modifier is specifically intended to be used on a target or
+      // target data directive to prevent data from being deallocated during the
+      // associated region.  It is not useful on a target enter data or target
+      // exit data directive because they are inherently dynamic not structured.
+      if ((DKind == OMPD_target_enter_data || DKind == OMPD_target_exit_data) &&
+          HasHoldModifier) {
+        SemaRef.Diag(StartLoc,
+                     diag::err_omp_invalid_map_type_modifier_for_directive)
+            << getOpenMPSimpleClauseTypeName(OMPC_map, OMPC_MAP_MODIFIER_hold)
+            << getOpenMPDirectiveName(DKind);
+        continue;
+      }
+
       // target, target data
       // OpenMP 5.0 [2.12.2, Restrictions, p. 163]
       // OpenMP 5.0 [2.12.5, Restrictions, p. 174]
@@ -17734,7 +17751,7 @@ OMPClause *Sema::ActOnOpenMPMapClause(
   OpenMPMapModifierKind Modifiers[] = {
       OMPC_MAP_MODIFIER_unknown, OMPC_MAP_MODIFIER_unknown,
       OMPC_MAP_MODIFIER_unknown, OMPC_MAP_MODIFIER_unknown,
-      OMPC_MAP_MODIFIER_unknown};
+      OMPC_MAP_MODIFIER_unknown, OMPC_MAP_MODIFIER_unknown};
   SourceLocation ModifiersLoc[NumberOfOMPMapClauseModifiers];
 
   // Process map-type-modifiers, flag errors for duplicate modifiers.

clang/test/OpenMP/target_enter_data_map_messages.c

@@ -25,5 +25,10 @@ int main(int argc, char **argv) {
   #pragma omp target enter data map(release: r) // expected-error {{map type 'release' is not allowed for '#pragma omp target enter data'}}
   #pragma omp target enter data map(delete: r) // expected-error {{map type 'delete' is not allowed for '#pragma omp target enter data'}}
 
+  // expected-error@+1 {{map type modifier 'hold' is not allowed for '#pragma omp target enter data'}}
+  #pragma omp target enter data map(hold, alloc: r)
+  // expected-error@+1 {{map type modifier 'hold' is not allowed for '#pragma omp target enter data'}}
+  #pragma omp target enter data map(hold, to: r)
+
   return 0;
 }

clang/test/OpenMP/target_exit_data_map_messages.c

@@ -18,5 +18,12 @@ int main(int argc, char **argv) {
   #pragma omp target exit data map(always, alloc: r) // expected-error {{map type 'alloc' is not allowed for '#pragma omp target exit data'}}
   #pragma omp target exit data map(to: r) // expected-error {{map type 'to' is not allowed for '#pragma omp target exit data'}}
 
+  // expected-error@+1 {{map type modifier 'hold' is not allowed for '#pragma omp target exit data'}}
+  #pragma omp target exit data map(hold, from: r)
+  // expected-error@+1 {{map type modifier 'hold' is not allowed for '#pragma omp target exit data'}}
+  #pragma omp target exit data map(hold, release: r)
+  // expected-error@+1 {{map type modifier 'hold' is not allowed for '#pragma omp target exit data'}}
+  #pragma omp target exit data map(hold, delete: r)
+
   return 0;
 }

openmp/libomptarget/include/omptarget.h

@@ -56,6 +56,12 @@ enum tgt_map_type {
   // TODO: Currently, it is not well tested outside of translations from OpenACC
   // to OpenMP, so it is not yet recommended for general use in OpenMP code.
   OMP_TGT_MAPTYPE_NO_ALLOC        = 0x2000,
+  // use a separate reference counter so that the data cannot be deallocated
+  // within the structured region
+  // This is an OpenMP extension for the sake of OpenACC support.
+  // TODO: Currently, it is not well tested outside of translations from OpenACC
+  // to OpenMP, so it is not yet recommended for general use in OpenMP code.
+  OMP_TGT_MAPTYPE_HOLD            = 0x4000,
   // member of struct, member given by [16 MSBs] - 1
   OMP_TGT_MAPTYPE_MEMBER_OF       = 0xffff000000000000
 };

openmp/libomptarget/src/api.cpp

@@ -112,7 +112,8 @@ EXTERN int omp_target_is_present(void *ptr, int device_num) {
   DeviceTy& Device = Devices[device_num];
   bool IsLast; // not used
   bool IsHostPtr;
-  void *TgtPtr = Device.getTgtPtrBegin(ptr, 0, IsLast, false, IsHostPtr);
+  void *TgtPtr = Device.getTgtPtrBegin(ptr, 0, IsLast, /*UpdateRefCount=*/false,
+                                       /*UseHoldRefCount=*/false, IsHostPtr);
   int rc = (TgtPtr != NULL);
   // Under unified memory the host pointer can be returned by the
   // getTgtPtrBegin() function which means that there is no device

openmp/libomptarget/src/device.cpp

@@ -47,6 +47,7 @@ int DeviceTy::associatePtr(void *HstPtrBegin, void *TgtPtrBegin, int64_t Size) {
                               (uintptr_t) HstPtrBegin /*HstPtrBegin*/,
                               (uintptr_t) HstPtrBegin + Size /*HstPtrEnd*/,
                               (uintptr_t) TgtPtrBegin /*TgtPtrBegin*/,
+                              false /*UseHoldRefCount*/,
                               true /*IsRefCountINF*/);
 
   DP("Creating new map entry: HstBase=" DPxMOD ", HstBegin=" DPxMOD ", HstEnd="
@@ -163,7 +164,8 @@ void *DeviceTy::getOrAllocTgtPtr(void *HstPtrBegin, void *HstPtrBase,
                                  int64_t Size, bool &IsNew, bool &IsHostPtr,
                                  bool IsImplicit, bool UpdateRefCount,
                                  bool HasCloseModifier, bool HasPresentModifier,
-                                 bool HasNoAllocModifier) {
+                                 bool HasNoAllocModifier,
+                                 bool HasHoldModifier) {
   void *rc = NULL;
   IsHostPtr = false;
   IsNew = false;
@@ -180,7 +182,7 @@ void *DeviceTy::getOrAllocTgtPtr(void *HstPtrBegin, void *HstPtrBase,
     IsNew = false;
 
     if (UpdateRefCount)
-      HT.incRefCount();
+      HT.incRefCount(HasHoldModifier);
 
     uintptr_t tp = HT.TgtPtrBegin + ((uintptr_t)HstPtrBegin - HT.HstPtrBegin);
     DP("Mapping exists%s with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD ", "
@@ -270,7 +272,8 @@ void *DeviceTy::getOrAllocTgtPtr(void *HstPtrBegin, void *HstPtrBase,
        DPxPTR((uintptr_t)HstPtrBegin + Size), DPxPTR(tp));
     HostDataToTargetMap.emplace(
         HostDataToTargetTy((uintptr_t)HstPtrBase, (uintptr_t)HstPtrBegin,
-                           (uintptr_t)HstPtrBegin + Size, tp));
+                           (uintptr_t)HstPtrBegin + Size, tp, HasHoldModifier,
+                           /*IsINF=*/false));
     rc = (void *)tp;
   }
 
@@ -282,8 +285,8 @@ void *DeviceTy::getOrAllocTgtPtr(void *HstPtrBegin, void *HstPtrBase,
 // Return the target pointer begin (where the data will be moved).
 // Decrement the reference counter if called from targetDataEnd.
 void *DeviceTy::getTgtPtrBegin(void *HstPtrBegin, int64_t Size, bool &IsLast,
-                               bool UpdateRefCount, bool &IsHostPtr,
-                               bool MustContain) {
+                               bool UpdateRefCount, bool UseHoldRefCount,
+                               bool &IsHostPtr, bool MustContain) {
   void *rc = NULL;
   IsHostPtr = false;
   IsLast = false;
@@ -293,10 +296,10 @@ void *DeviceTy::getTgtPtrBegin(void *HstPtrBegin, int64_t Size, bool &IsLast,
   if (lr.Flags.IsContained ||
       (!MustContain && (lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter))) {
     auto &HT = *lr.Entry;
-    IsLast = HT.getRefCount() == 1;
+    IsLast = HT.getRefCount() == 1 && HT.getRefCount(UseHoldRefCount) == 1;
 
-    if (!IsLast && UpdateRefCount)
-      HT.decRefCount();
+    if (!IsLast && UpdateRefCount && HT.getRefCount(UseHoldRefCount) > 0)
+      HT.decRefCount(UseHoldRefCount);
 
     uintptr_t tp = HT.TgtPtrBegin + ((uintptr_t)HstPtrBegin - HT.HstPtrBegin);
     DP("Mapping exists with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD ", "
@@ -333,7 +336,7 @@ void *DeviceTy::getTgtPtrBegin(void *HstPtrBegin, int64_t Size) {
 }
 
 int DeviceTy::deallocTgtPtr(void *HstPtrBegin, int64_t Size, bool ForceDelete,
-                            bool HasCloseModifier) {
+                            bool HasCloseModifier, bool HasHoldModifier) {
   if (RTLs->RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY && !HasCloseModifier)
     return OFFLOAD_SUCCESS;
   // Check if the pointer is contained in any sub-nodes.
@@ -344,7 +347,7 @@ int DeviceTy::deallocTgtPtr(void *HstPtrBegin, int64_t Size, bool ForceDelete,
     auto &HT = *lr.Entry;
     if (ForceDelete)
       HT.resetRefCount();
-    if (HT.decRefCount() == 0) {
+    if (HT.decRefCount(HasHoldModifier) == 0) {
       DP("Deleting tgt data " DPxMOD " of size %" PRId64 "\n",
           DPxPTR(HT.TgtPtrBegin), Size);
 #if OMPT_SUPPORT

openmp/libomptarget/src/device.h

@@ -40,44 +40,69 @@ struct HostDataToTargetTy {
 
 private:
   /// use mutable to allow modification via std::set iterator which is const.
+  ///@{
   mutable uint64_t RefCount;
+  mutable uint64_t HoldRefCount;
+  ///@}
   static const uint64_t INFRefCount = ~(uint64_t)0;
 
 public:
   HostDataToTargetTy(uintptr_t BP, uintptr_t B, uintptr_t E, uintptr_t TB,
-      bool IsINF = false)
-      : HstPtrBase(BP), HstPtrBegin(B), HstPtrEnd(E),
-        TgtPtrBegin(TB), RefCount(IsINF ? INFRefCount : 1) {}
+                     bool UseHoldRefCount, bool IsINF)
+      : HstPtrBase(BP), HstPtrBegin(B), HstPtrEnd(E), TgtPtrBegin(TB),
+        RefCount(IsINF ? INFRefCount : !UseHoldRefCount),
+        HoldRefCount(UseHoldRefCount) {}
 
+  /// Get the total reference count.
   uint64_t getRefCount() const {
-    return RefCount;
+    if (RefCount == INFRefCount)
+      return RefCount;
+    return RefCount + HoldRefCount;
   }
 
+  /// Get a specific reference count.
+  uint64_t getRefCount(bool UseHoldRefCount) const {
+    return UseHoldRefCount ? HoldRefCount : RefCount;
+  }
+
+  /// Reset the dynamic reference count only and return the total reference
+  /// count.
   uint64_t resetRefCount() const {
     if (RefCount != INFRefCount)
       RefCount = 1;
 
-    return RefCount;
+    return getRefCount();
   }
 
-  uint64_t incRefCount() const {
-    if (RefCount != INFRefCount) {
+  /// Increment the specified reference count and return the total reference
+  /// count.
+  uint64_t incRefCount(bool UseHoldRefCount) const {
+    if (UseHoldRefCount) {
+      ++HoldRefCount;
+      assert(HoldRefCount < INFRefCount && "hold refcount overflow");
+    } else if (RefCount != INFRefCount) {
       ++RefCount;
       assert(RefCount < INFRefCount && "refcount overflow");
     }
 
-    return RefCount;
+    return getRefCount();
   }
 
-  uint64_t decRefCount() const {
-    if (RefCount != INFRefCount) {
+  /// Decrement the specified reference count and return the total reference
+  /// count.
+  uint64_t decRefCount(bool UseHoldRefCount) const {
+    if (UseHoldRefCount) {
+      assert(HoldRefCount > 0 && "hold refcount underflow");
+      --HoldRefCount;
+    } else if (RefCount != INFRefCount) {
       assert(RefCount > 0 && "refcount underflow");
       --RefCount;
     }
 
-    return RefCount;
+    return getRefCount();
   }
 
+  /// Is the dynamic (and thus total) reference count infinite?
   bool isRefCountInf() const {
     return RefCount == INFRefCount;
   }
@@ -196,13 +221,14 @@ struct DeviceTy {
   void *getOrAllocTgtPtr(void *HstPtrBegin, void *HstPtrBase, int64_t Size,
                          bool &IsNew, bool &IsHostPtr, bool IsImplicit,
                          bool UpdateRefCount, bool HasCloseModifier,
-                         bool HasPresentModifier, bool HasNoAllocModifier);
+                         bool HasPresentModifier, bool HasNoAllocModifier,
+                         bool HasHoldModifier);
   void *getTgtPtrBegin(void *HstPtrBegin, int64_t Size);
   void *getTgtPtrBegin(void *HstPtrBegin, int64_t Size, bool &IsLast,
-                       bool UpdateRefCount, bool &IsHostPtr,
-                       bool MustContain = false);
+                       bool UpdateRefCount, bool UseHoldRefCount,
+                       bool &IsHostPtr, bool MustContain = false);
   int deallocTgtPtr(void *TgtPtrBegin, int64_t Size, bool ForceDelete,
-                    bool HasCloseModifier = false);
+                    bool HasCloseModifier, bool HasHoldModifier);
   int associatePtr(void *HstPtrBegin, void *TgtPtrBegin, int64_t Size);
   int disassociatePtr(void *HstPtrBegin);