diff --git a/openmp/libomptarget/src/device.cpp b/openmp/libomptarget/src/device.cpp index c70d9e89757261..62d694ec8529e0 100644 --- a/openmp/libomptarget/src/device.cpp +++ b/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(); diff --git a/openmp/libomptarget/src/device.h b/openmp/libomptarget/src/device.h index ea87a4b270f0d6..16120537c1f61f 100644 --- a/openmp/libomptarget/src/device.h +++ b/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, diff --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp index fe50563d94a6c7..55b840a611fed2 100644 --- a/openmp/libomptarget/src/omptarget.cpp +++ b/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 diff --git a/openmp/libomptarget/test/unified_shared_memory/associate_ptr.c b/openmp/libomptarget/test/unified_shared_memory/associate_ptr.c index 7911046f5f3b12..c25c557b3b6d37 100644 --- a/openmp/libomptarget/test/unified_shared_memory/associate_ptr.c +++ b/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));