[OpenMP] Use IsHostPtr where needed for targetDataBegin

As discussed in D105990, without this patch, `targetDataBegin`
determines whether to transfer data (as opposed to assuming it's in
shared memory) using the condition `!UseUSM || HasCloseModifier`.
However, this condition is broken if use of discrete memory was forced
by `omp_target_associate_ptr`.  This patch extends
`unified_shared_memory/associate_ptr.c` to reveal this case, and it
fixes it using `!IsHostPtr` in `DeviceTy::getTargetPointer` to replace
this condition.

Reviewed By: grokos

Differential Revision: https://reviews.llvm.org/D107927

openmp/libomptarget/src/device.cpp

@@ -177,10 +177,11 @@ LookupResult DeviceTy::lookupMapping(void *HstPtrBegin, int64_t Size) {
 
 TargetPointerResultTy
 DeviceTy::getTargetPointer(void *HstPtrBegin, void *HstPtrBase, int64_t Size,
-                           map_var_info_t HstPtrName, MoveDataStateTy MoveData,
-                           bool IsImplicit, bool UpdateRefCount,
-                           bool HasCloseModifier, bool HasPresentModifier,
-                           bool HasHoldModifier, AsyncInfoTy &AsyncInfo) {
+                           map_var_info_t HstPtrName, bool HasFlagTo,
+                           bool HasFlagAlways, bool IsImplicit,
+                           bool UpdateRefCount, bool HasCloseModifier,
+                           bool HasPresentModifier, bool HasHoldModifier,
+                           AsyncInfoTy &AsyncInfo) {
   void *TargetPointer = nullptr;
   bool IsHostPtr = false;
   bool IsNew = false;
@@ -272,12 +273,9 @@ DeviceTy::getTargetPointer(void *HstPtrBegin, void *HstPtrBase, int64_t Size,
     TargetPointer = (void *)Ptr;
   }
 
-  if (IsNew && MoveData == MoveDataStateTy::UNKNOWN)
-    MoveData = MoveDataStateTy::REQUIRED;
-
   // If the target pointer is valid, and we need to transfer data, issue the
   // data transfer.
-  if (TargetPointer && (MoveData == MoveDataStateTy::REQUIRED)) {
+  if (TargetPointer && !IsHostPtr && HasFlagTo && (IsNew || HasFlagAlways)) {
     // Lock the entry before releasing the mapping table lock such that another
     // thread that could issue data movement will get the right result.
     Entry->lock();

openmp/libomptarget/src/device.h

@@ -226,8 +226,6 @@ struct PendingCtorDtorListsTy {
 typedef std::map<__tgt_bin_desc *, PendingCtorDtorListsTy>
     PendingCtorsDtorsPerLibrary;
 
-enum class MoveDataStateTy : uint32_t { REQUIRED, NONE, UNKNOWN };
-
 struct DeviceTy {
   int32_t DeviceID;
   RTLInfoTy *RTL;
@@ -264,20 +262,20 @@ struct DeviceTy {
   LookupResult lookupMapping(void *HstPtrBegin, int64_t Size);
   /// Get the target pointer based on host pointer begin and base. If the
   /// mapping already exists, the target pointer will be returned directly. In
-  /// addition, if \p MoveData is true, the memory region pointed by \p
-  /// HstPtrBegin of size \p Size will also be transferred to the device. If the
-  /// mapping doesn't exist, and if unified memory is not enabled, a new mapping
-  /// will be created and the data will also be transferred accordingly. nullptr
-  /// will be returned because of any of following reasons:
+  /// addition, if required, the memory region pointed by \p HstPtrBegin of size
+  /// \p Size will also be transferred to the device. If the mapping doesn't
+  /// exist, and if unified shared memory is not enabled, a new mapping will be
+  /// created and the data will also be transferred accordingly. nullptr will be
+  /// returned because of any of following reasons:
   /// - Data allocation failed;
   /// - The user tried to do an illegal mapping;
   /// - Data transfer issue fails.
   TargetPointerResultTy
   getTargetPointer(void *HstPtrBegin, void *HstPtrBase, int64_t Size,
-                   map_var_info_t HstPtrName, MoveDataStateTy MoveData,
-                   bool IsImplicit, bool UpdateRefCount, bool HasCloseModifier,
-                   bool HasPresentModifier, bool HasHoldModifier,
-                   AsyncInfoTy &AsyncInfo);
+                   map_var_info_t HstPtrName, bool HasFlagTo,
+                   bool HasFlagAlways, bool IsImplicit, bool UpdateRefCount,
+                   bool HasCloseModifier, bool HasPresentModifier,
+                   bool HasHoldModifier, AsyncInfoTy &AsyncInfo);
   void *getTgtPtrBegin(void *HstPtrBegin, int64_t Size);
   void *getTgtPtrBegin(void *HstPtrBegin, int64_t Size, bool &IsLast,
                        bool UpdateRefCount, bool UseHoldRefCount,

openmp/libomptarget/src/omptarget.cpp

@@ -490,9 +490,9 @@ int targetDataBegin(ident_t *loc, DeviceTy &Device, int32_t arg_num,
       // PTR_AND_OBJ entry is handled below, and so the allocation might fail
       // when HasPresentModifier.
       Pointer_TPR = Device.getTargetPointer(
-          HstPtrBase, HstPtrBase, sizeof(void *), nullptr,
-          MoveDataStateTy::NONE, IsImplicit, UpdateRef, HasCloseModifier,
-          HasPresentModifier, HasHoldModifier, AsyncInfo);
+          HstPtrBase, HstPtrBase, sizeof(void *), /*HstPtrName=*/nullptr,
+          /*HasFlagTo=*/false, /*HasFlagAlways=*/false, IsImplicit, UpdateRef,
+          HasCloseModifier, HasPresentModifier, HasHoldModifier, AsyncInfo);
       PointerTgtPtrBegin = Pointer_TPR.TargetPointer;
       IsHostPtr = Pointer_TPR.Flags.IsHostPointer;
       if (!PointerTgtPtrBegin) {
@@ -514,18 +514,13 @@ int targetDataBegin(ident_t *loc, DeviceTy &Device, int32_t arg_num,
           (!FromMapper || i != 0); // subsequently update ref count of pointee
     }
 
-    MoveDataStateTy MoveData = MoveDataStateTy::NONE;
-    const bool UseUSM = PM->RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY;
     const bool HasFlagTo = arg_types[i] & OMP_TGT_MAPTYPE_TO;
     const bool HasFlagAlways = arg_types[i] & OMP_TGT_MAPTYPE_ALWAYS;
-    if (HasFlagTo && (!UseUSM || HasCloseModifier))
-      MoveData = HasFlagAlways ? MoveDataStateTy::REQUIRED
-                               : MoveDataStateTy::UNKNOWN;
-
-    auto TPR = Device.getTargetPointer(
-        HstPtrBegin, HstPtrBase, data_size, HstPtrName, MoveData, IsImplicit,
-        UpdateRef, HasCloseModifier, HasPresentModifier, HasHoldModifier,
-        AsyncInfo);
+    auto TPR = Device.getTargetPointer(HstPtrBegin, HstPtrBase, data_size,
+                                       HstPtrName, HasFlagTo, HasFlagAlways,
+                                       IsImplicit, UpdateRef, HasCloseModifier,
+                                       HasPresentModifier, HasHoldModifier,
+                                       AsyncInfo);
     void *TgtPtrBegin = TPR.TargetPointer;
     IsHostPtr = TPR.Flags.IsHostPointer;
     // If data_size==0, then the argument could be a zero-length pointer to

openmp/libomptarget/test/unified_shared_memory/associate_ptr.c

@@ -25,9 +25,9 @@ int main(int argc, char *argv[]) {
   // specified.  It must check whether x was previously placed in device memory
   // by, for example, omp_target_associate_ptr.
   #pragma omp target map(always, tofrom: x)
-  x = 20;
+  x += 1;
 
-  // CHECK: x=20
+  // CHECK: x=11
   printf("x=%d\n", x);
   // CHECK: present: 1
   printf("present: %d\n", omp_target_is_present(&x, dev));