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;
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();
typedef std::map<__tgt_bin_desc *, PendingCtorDtorListsTy>
PendingCtorsDtorsPerLibrary;
-enum class MoveDataStateTy : uint32_t { REQUIRED, NONE, UNKNOWN };
-
struct DeviceTy {
int32_t DeviceID;
RTLInfoTy *RTL;
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,
// 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) {
(!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
// 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));