intel/llvm
1
0
Fork 0
mirror of https://github.com/intel/llvm.git synced 2026-01-16 13:35:38 +08:00
Code Issues Packages Projects Releases Wiki Activity

[OpenMP][OpenACC] Implement ompx_hold map type modifier extension in runtime (2/2)

Browse Source 
This patch implements OpenMP runtime support for an original OpenMP
extension we have developed to support OpenACC: the `ompx_hold` map
type modifier.  The previous patch in this series, D106509, implements
Clang support and documents the new functionality in detail.

Reviewed By: grokos

Differential Revision: https://reviews.llvm.org/D106510
This commit is contained in:
Joel E. Denny
2021-08-31 15:21:16 -04:00
parent 83ddfa0d22
commit ec1ebcd302
11 changed files with 877 additions and 110 deletions
Download Patch File Download Diff File Expand all files Collapse all files

4
openmp/libomptarget/include/omptarget.h
View File

@@ -56,6 +56,10 @@ enum tgt_map_type {
OMP_TGT_MAPTYPE_CLOSE = 0x400,
// runtime error if not already allocated
OMP_TGT_MAPTYPE_PRESENT = 0x1000,
// use a separate reference counter so that the data cannot be unmapped within
// the structured region
// This is an OpenMP extension for the sake of OpenACC support.
OMP_TGT_MAPTYPE_OMPX_HOLD = 0x2000,
// descriptor for non-contiguous target-update
OMP_TGT_MAPTYPE_NON_CONTIG = 0x100000000000,
// member of struct, member given by [16 MSBs] - 1

3
openmp/libomptarget/src/api.cpp
View File

@@ -106,7 +106,8 @@ EXTERN int omp_target_is_present(const void *ptr, int device_num) {
bool IsLast; // not used
bool IsHostPtr;
void *TgtPtr = Device.getTgtPtrBegin(const_cast<void *>(ptr), 0, IsLast,
false, IsHostPtr);
/*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

95
openmp/libomptarget/src/device.cpp
View File

@@ -82,14 +82,16 @@ int DeviceTy::associatePtr(void *HstPtrBegin, void *TgtPtrBegin, int64_t Size) {
/*HstPtrBase=*/(uintptr_t)HstPtrBegin,
/*HstPtrBegin=*/(uintptr_t)HstPtrBegin,
/*HstPtrEnd=*/(uintptr_t)HstPtrBegin + Size,
/*TgtPtrBegin=*/(uintptr_t)TgtPtrBegin, /*Name=*/nullptr,
/*TgtPtrBegin=*/(uintptr_t)TgtPtrBegin,
/*UseHoldRefCount=*/false, /*Name=*/nullptr,
/*IsRefCountINF=*/true)
.first;
DP("Creating new map entry: HstBase=" DPxMOD ", HstBegin=" DPxMOD
", HstEnd=" DPxMOD ", TgtBegin=" DPxMOD ", RefCount=%s\n",
", HstEnd=" DPxMOD ", TgtBegin=" DPxMOD ", DynRefCount=%s, "
"HoldRefCount=%s\n",
DPxPTR(newEntry.HstPtrBase), DPxPTR(newEntry.HstPtrBegin),
DPxPTR(newEntry.HstPtrEnd), DPxPTR(newEntry.TgtPtrBegin),
newEntry.refCountToStr().c_str());
newEntry.dynRefCountToStr().c_str(), newEntry.holdRefCountToStr().c_str());
(void)newEntry;
DataMapMtx.unlock();
@@ -103,7 +105,13 @@ int DeviceTy::disassociatePtr(void *HstPtrBegin) {
auto search = HostDataToTargetMap.find(HstPtrBeginTy{(uintptr_t)HstPtrBegin});
if (search != HostDataToTargetMap.end()) {
// Mapping exists
if (search->isRefCountInf()) {
if (search->getHoldRefCount()) {
// This is based on OpenACC 3.1, sec 3.2.33 "acc_unmap_data", L3656-3657:
// "It is an error to call acc_unmap_data if the structured reference
// count for the pointer is not zero."
REPORT("Trying to disassociate a pointer with a non-zero hold reference "
"count\n");
} else if (search->isDynRefCountInf()) {
DP("Association found, removing it\n");
HostDataToTargetMap.erase(search);
DataMapMtx.unlock();
@@ -112,11 +120,12 @@ int DeviceTy::disassociatePtr(void *HstPtrBegin) {
REPORT("Trying to disassociate a pointer which was not mapped via "
"omp_target_associate_ptr\n");
}
} else {
REPORT("Association not found\n");
}
// Mapping not found
DataMapMtx.unlock();
REPORT("Association not found\n");
return OFFLOAD_FAIL;
}
@@ -171,7 +180,7 @@ DeviceTy::getTargetPointer(void *HstPtrBegin, void *HstPtrBase, int64_t Size,
map_var_info_t HstPtrName, MoveDataStateTy MoveData,
bool IsImplicit, bool UpdateRefCount,
bool HasCloseModifier, bool HasPresentModifier,
AsyncInfoTy &AsyncInfo) {
bool HasHoldModifier, AsyncInfoTy &AsyncInfo) {
void *TargetPointer = nullptr;
bool IsHostPtr = false;
bool IsNew = false;
@@ -188,21 +197,26 @@ DeviceTy::getTargetPointer(void *HstPtrBegin, void *HstPtrBase, int64_t Size,
if (LR.Flags.IsContained ||
((LR.Flags.ExtendsBefore || LR.Flags.ExtendsAfter) && IsImplicit)) {
auto &HT = *LR.Entry;
assert(HT.getRefCount() > 0 && "expected existing RefCount > 0");
if (UpdateRefCount)
const char *RefCountAction;
assert(HT.getTotalRefCount() > 0 && "expected existing RefCount > 0");
if (UpdateRefCount) {
// After this, RefCount > 1.
HT.incRefCount();
else
HT.incRefCount(HasHoldModifier);
RefCountAction = " (incremented)";
} else {
// It might have been allocated with the parent, but it's still new.
IsNew = HT.getRefCount() == 1;
IsNew = HT.getTotalRefCount() == 1;
RefCountAction = " (update suppressed)";
}
const char *DynRefCountAction = HasHoldModifier ? "" : RefCountAction;
const char *HoldRefCountAction = HasHoldModifier ? RefCountAction : "";
uintptr_t Ptr = HT.TgtPtrBegin + ((uintptr_t)HstPtrBegin - HT.HstPtrBegin);
INFO(OMP_INFOTYPE_MAPPING_EXISTS, DeviceID,
"Mapping exists%s with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD
", "
"Size=%" PRId64 ", RefCount=%s (%s), Name=%s\n",
", Size=%" PRId64 ", DynRefCount=%s%s, HoldRefCount=%s%s, Name=%s\n",
(IsImplicit ? " (implicit)" : ""), DPxPTR(HstPtrBegin), DPxPTR(Ptr),
Size, HT.refCountToStr().c_str(),
UpdateRefCount ? "incremented" : "update suppressed",
Size, HT.dynRefCountToStr().c_str(), DynRefCountAction,
HT.holdRefCountToStr().c_str(), HoldRefCountAction,
(HstPtrName) ? getNameFromMapping(HstPtrName).c_str() : "unknown");
TargetPointer = (void *)Ptr;
} else if ((LR.Flags.ExtendsBefore || LR.Flags.ExtendsAfter) && !IsImplicit) {
@@ -245,13 +259,15 @@ DeviceTy::getTargetPointer(void *HstPtrBegin, void *HstPtrBase, int64_t Size,
uintptr_t Ptr = (uintptr_t)allocData(Size, HstPtrBegin);
Entry = HostDataToTargetMap
.emplace((uintptr_t)HstPtrBase, (uintptr_t)HstPtrBegin,
(uintptr_t)HstPtrBegin + Size, Ptr, HstPtrName)
(uintptr_t)HstPtrBegin + Size, Ptr, HasHoldModifier,
HstPtrName)
.first;
INFO(OMP_INFOTYPE_MAPPING_CHANGED, DeviceID,
"Creating new map entry with "
"HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD ", Size=%ld, "
"RefCount=%s, Name=%s\n",
DPxPTR(HstPtrBegin), DPxPTR(Ptr), Size, Entry->refCountToStr().c_str(),
"DynRefCount=%s, HoldRefCount=%s, Name=%s\n",
DPxPTR(HstPtrBegin), DPxPTR(Ptr), Size,
Entry->dynRefCountToStr().c_str(), Entry->holdRefCountToStr().c_str(),
(HstPtrName) ? getNameFromMapping(HstPtrName).c_str() : "unknown");
TargetPointer = (void *)Ptr;
}
@@ -295,8 +311,9 @@ DeviceTy::getTargetPointer(void *HstPtrBegin, void *HstPtrBase, int64_t Size,
// 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 ForceDelete) {
bool UpdateRefCount, bool UseHoldRefCount,
bool &IsHostPtr, bool MustContain,
bool ForceDelete) {
void *rc = NULL;
IsHostPtr = false;
IsLast = false;
@@ -306,35 +323,39 @@ 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;
// We do not decrement the reference count to zero here. deallocTgtPtr does
// that atomically with removing the mapping. Otherwise, before this thread
// We do not zero the total reference count here. deallocTgtPtr does that
// atomically with removing the mapping. Otherwise, before this thread
// removed the mapping in deallocTgtPtr, another thread could retrieve the
// mapping, increment and decrement back to zero, and then both threads
// would try to remove the mapping, resulting in a double free.
IsLast = HT.decShouldRemove(ForceDelete);
IsLast = HT.decShouldRemove(UseHoldRefCount, ForceDelete);
const char *RefCountAction;
if (!UpdateRefCount) {
RefCountAction = "update suppressed";
RefCountAction = " (update suppressed)";
} else if (ForceDelete) {
HT.resetRefCount();
assert(IsLast == HT.decShouldRemove() &&
HT.resetRefCount(UseHoldRefCount);
assert(IsLast == HT.decShouldRemove(UseHoldRefCount) &&
"expected correct IsLast prediction for reset");
if (IsLast)
RefCountAction = "reset, deferred final decrement";
else
RefCountAction = "reset";
RefCountAction = " (reset, deferred final decrement)";
else {
HT.decRefCount(UseHoldRefCount);
RefCountAction = " (reset)";
}
} else if (IsLast) {
RefCountAction = "deferred final decrement";
RefCountAction = " (deferred final decrement)";
} else {
RefCountAction = "decremented";
HT.decRefCount();
HT.decRefCount(UseHoldRefCount);
RefCountAction = " (decremented)";
}
const char *DynRefCountAction = UseHoldRefCount ? "" : RefCountAction;
const char *HoldRefCountAction = UseHoldRefCount ? RefCountAction : "";
uintptr_t tp = HT.TgtPtrBegin + ((uintptr_t)HstPtrBegin - HT.HstPtrBegin);
INFO(OMP_INFOTYPE_MAPPING_EXISTS, DeviceID,
"Mapping exists with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD ", "
"Size=%" PRId64 ", RefCount=%s (%s)\n",
DPxPTR(HstPtrBegin), DPxPTR(tp), Size, HT.refCountToStr().c_str(),
RefCountAction);
"Size=%" PRId64 ", DynRefCount=%s%s, HoldRefCount=%s%s\n",
DPxPTR(HstPtrBegin), DPxPTR(tp), Size, HT.dynRefCountToStr().c_str(),
DynRefCountAction, HT.holdRefCountToStr().c_str(), HoldRefCountAction);
rc = (void *)tp;
} else if (PM->RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY) {
// If the value isn't found in the mapping and unified shared memory
@@ -366,7 +387,7 @@ void *DeviceTy::getTgtPtrBegin(void *HstPtrBegin, int64_t Size) {
}
int DeviceTy::deallocTgtPtr(void *HstPtrBegin, int64_t Size,
bool HasCloseModifier) {
bool HasCloseModifier, bool HasHoldModifier) {
if (PM->RTLs.RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY &&
!HasCloseModifier)
return OFFLOAD_SUCCESS;
@@ -376,7 +397,7 @@ int DeviceTy::deallocTgtPtr(void *HstPtrBegin, int64_t Size,
LookupResult lr = lookupMapping(HstPtrBegin, Size);
if (lr.Flags.IsContained || lr.Flags.ExtendsBefore || lr.Flags.ExtendsAfter) {
auto &HT = *lr.Entry;
if (HT.decRefCount() == 0) {
if (HT.decRefCount(HasHoldModifier) == 0) {
DP("Deleting tgt data " DPxMOD " of size %" PRId64 "\n",
DPxPTR(HT.TgtPtrBegin), Size);
deleteData((void *)HT.TgtPtrBegin);

129
openmp/libomptarget/src/device.h
View File

@@ -50,9 +50,30 @@ struct HostDataToTargetTy {
uintptr_t TgtPtrBegin; // target info.
private:
/// use mutable to allow modification via std::set iterator which is const.
mutable uint64_t RefCount;
/// The dynamic reference count is the standard reference count as of OpenMP
/// 4.5. The hold reference count is an OpenMP extension for the sake of
/// OpenACC support.
///
/// The 'ompx_hold' map type modifier is permitted only on "omp target" and
/// "omp target data", and "delete" is permitted only on "omp target exit
/// data" and associated runtime library routines. As a result, we really
/// need to implement "reset" functionality only for the dynamic reference
/// counter. Likewise, only the dynamic reference count can be infinite
/// because, for example, omp_target_associate_ptr and "omp declare target
/// link" operate only on it. Nevertheless, it's actually easier to follow
/// the code (and requires less assertions for special cases) when we just
/// implement these features generally across both reference counters here.
/// Thus, it's the users of this class that impose those restrictions.
///
/// Use mutable to allow modification via std::set iterator which is const.
///@{
mutable uint64_t DynRefCount;
mutable uint64_t HoldRefCount;
///@}
static const uint64_t INFRefCount = ~(uint64_t)0;
static std::string refCountToStr(uint64_t RefCount) {
return RefCount == INFRefCount ? "INF" : std::to_string(RefCount);
}
/// This mutex will be locked when data movement is issued. For targets that
/// doesn't support async data movement, this mutex can guarantee that after
/// it is released, memory region on the target is update to date. For targets
@@ -63,50 +84,82 @@ private:
public:
HostDataToTargetTy(uintptr_t BP, uintptr_t B, uintptr_t E, uintptr_t TB,
map_var_info_t Name = nullptr, bool IsINF = false)
bool UseHoldRefCount, map_var_info_t Name = nullptr,
bool IsINF = false)
: HstPtrBase(BP), HstPtrBegin(B), HstPtrEnd(E), HstPtrName(Name),
TgtPtrBegin(TB), RefCount(IsINF ? INFRefCount : 1),
TgtPtrBegin(TB), DynRefCount(UseHoldRefCount ? 0
: IsINF ? INFRefCount
: 1),
HoldRefCount(!UseHoldRefCount ? 0
: IsINF ? INFRefCount
: 1),
UpdateMtx(std::make_shared<std::mutex>()) {}
uint64_t getRefCount() const { return RefCount; }
uint64_t resetRefCount() const {
if (RefCount != INFRefCount)
RefCount = 1;
return RefCount;
/// Get the total reference count. This is smarter than just getDynRefCount()
/// + getHoldRefCount() because it handles the case where at least one is
/// infinity and the other is non-zero.
uint64_t getTotalRefCount() const {
if (DynRefCount == INFRefCount || HoldRefCount == INFRefCount)
return INFRefCount;
return DynRefCount + HoldRefCount;
}
uint64_t incRefCount() const {
if (RefCount != INFRefCount) {
++RefCount;
assert(RefCount < INFRefCount && "refcount overflow");
/// Get the dynamic reference count.
uint64_t getDynRefCount() const { return DynRefCount; }
/// Get the hold reference count.
uint64_t getHoldRefCount() const { return HoldRefCount; }
/// Reset the specified reference count unless it's infinity. Reset to 1
/// (even if currently 0) so it can be followed by a decrement.
void resetRefCount(bool UseHoldRefCount) const {
uint64_t &ThisRefCount = UseHoldRefCount ? HoldRefCount : DynRefCount;
if (ThisRefCount != INFRefCount)
ThisRefCount = 1;
}
/// Increment the specified reference count unless it's infinity.
void incRefCount(bool UseHoldRefCount) const {
uint64_t &ThisRefCount = UseHoldRefCount ? HoldRefCount : DynRefCount;
if (ThisRefCount != INFRefCount) {
++ThisRefCount;
assert(ThisRefCount < INFRefCount && "refcount overflow");
}
return RefCount;
}
uint64_t decRefCount() const {
if (RefCount != INFRefCount) {
assert(RefCount > 0 && "refcount underflow");
--RefCount;
/// Decrement the specified reference count unless it's infinity or zero, and
/// return the total reference count.
uint64_t decRefCount(bool UseHoldRefCount) const {
uint64_t &ThisRefCount = UseHoldRefCount ? HoldRefCount : DynRefCount;
uint64_t OtherRefCount = UseHoldRefCount ? DynRefCount : HoldRefCount;
if (ThisRefCount != INFRefCount) {
if (ThisRefCount > 0)
--ThisRefCount;
else
assert(OtherRefCount > 0 && "total refcount underflow");
}
return RefCount;
return getTotalRefCount();
}
bool isRefCountInf() const { return RefCount == INFRefCount; }
/// Is the dynamic (and thus the total) reference count infinite?
bool isDynRefCountInf() const { return DynRefCount == INFRefCount; }
std::string refCountToStr() const {
return isRefCountInf() ? "INF" : std::to_string(getRefCount());
}
/// Convert the dynamic reference count to a debug string.
std::string dynRefCountToStr() const { return refCountToStr(DynRefCount); }
/// Should one decrement of the reference count (after resetting it if
/// \c AfterReset) remove this mapping?
bool decShouldRemove(bool AfterReset = false) const {
/// Convert the hold reference count to a debug string.
std::string holdRefCountToStr() const { return refCountToStr(HoldRefCount); }
/// Should one decrement of the specified reference count (after resetting it
/// if \c AfterReset) remove this mapping?
bool decShouldRemove(bool UseHoldRefCount, bool AfterReset = false) const {
uint64_t ThisRefCount = UseHoldRefCount ? HoldRefCount : DynRefCount;
uint64_t OtherRefCount = UseHoldRefCount ? DynRefCount : HoldRefCount;
if (OtherRefCount > 0)
return false;
if (AfterReset)
return !isRefCountInf();
return getRefCount() == 1;
return ThisRefCount != INFRefCount;
return ThisRefCount == 1;
}
void lock() const { UpdateMtx->lock(); }
@@ -223,13 +276,15 @@ struct DeviceTy {
getTargetPointer(void *HstPtrBegin, void *HstPtrBase, int64_t Size,
map_var_info_t HstPtrName, MoveDataStateTy MoveData,
bool IsImplicit, bool UpdateRefCount, bool HasCloseModifier,
bool HasPresentModifier, AsyncInfoTy &AsyncInfo);
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 &IsHostPtr,
bool MustContain = false, bool ForceDelete = false);
int deallocTgtPtr(void *TgtPtrBegin, int64_t Size,
bool HasCloseModifier = false);
bool UpdateRefCount, bool UseHoldRefCount,
bool &IsHostPtr, bool MustContain = false,
bool ForceDelete = false);
int deallocTgtPtr(void *TgtPtrBegin, int64_t Size, bool HasCloseModifier,
bool HasHoldModifier);
int associatePtr(void *HstPtrBegin, void *TgtPtrBegin, int64_t Size);
int disassociatePtr(void *HstPtrBegin);

43
openmp/libomptarget/src/omptarget.cpp
View File

@@ -157,7 +157,8 @@ static int InitLibrary(DeviceTy &Device) {
(uintptr_t)CurrHostEntry->addr /*HstPtrBase*/,
(uintptr_t)CurrHostEntry->addr /*HstPtrBegin*/,
(uintptr_t)CurrHostEntry->addr + CurrHostEntry->size /*HstPtrEnd*/,
(uintptr_t)CurrDeviceEntry->addr /*TgtPtrBegin*/, nullptr,
(uintptr_t)CurrDeviceEntry->addr /*TgtPtrBegin*/,
false /*UseHoldRefCount*/, nullptr /*Name*/,
true /*IsRefCountINF*/);
}
}
@@ -465,6 +466,7 @@ int targetDataBegin(ident_t *loc, DeviceTy &Device, int32_t arg_num,
// a close map modifier was associated with a map that contained a to.
bool HasCloseModifier = arg_types[i] & OMP_TGT_MAPTYPE_CLOSE;
bool HasPresentModifier = arg_types[i] & OMP_TGT_MAPTYPE_PRESENT;
bool HasHoldModifier = arg_types[i] & OMP_TGT_MAPTYPE_OMPX_HOLD;
// UpdateRef is based on MEMBER_OF instead of TARGET_PARAM because if we
// have reached this point via __tgt_target_data_begin and not __tgt_target
// then no argument is marked as TARGET_PARAM ("omp target data map" is not
@@ -490,7 +492,7 @@ int targetDataBegin(ident_t *loc, DeviceTy &Device, int32_t arg_num,
Pointer_TPR = Device.getTargetPointer(
HstPtrBase, HstPtrBase, sizeof(void *), nullptr,
MoveDataStateTy::NONE, IsImplicit, UpdateRef, HasCloseModifier,
HasPresentModifier, AsyncInfo);
HasPresentModifier, HasHoldModifier, AsyncInfo);
PointerTgtPtrBegin = Pointer_TPR.TargetPointer;
IsHostPtr = Pointer_TPR.Flags.IsHostPointer;
if (!PointerTgtPtrBegin) {
@@ -522,7 +524,8 @@ int targetDataBegin(ident_t *loc, DeviceTy &Device, int32_t arg_num,
auto TPR = Device.getTargetPointer(
HstPtrBegin, HstPtrBase, data_size, HstPtrName, MoveData, IsImplicit,
UpdateRef, HasCloseModifier, HasPresentModifier, AsyncInfo);
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
@@ -608,10 +611,13 @@ struct DeallocTgtPtrInfo {
int64_t DataSize;
/// Whether it has \p close modifier
bool HasCloseModifier;
/// Whether it has \p ompx_hold modifier
bool HasHoldModifier;
DeallocTgtPtrInfo(void *HstPtr, int64_t Size, bool HasCloseModifier)
: HstPtrBegin(HstPtr), DataSize(Size),
HasCloseModifier(HasCloseModifier) {}
DeallocTgtPtrInfo(void *HstPtr, int64_t Size, bool HasCloseModifier,
bool HasHoldModifier)
: HstPtrBegin(HstPtr), DataSize(Size), HasCloseModifier(HasCloseModifier),
HasHoldModifier(HasHoldModifier) {}
};
} // namespace
@@ -678,11 +684,12 @@ int targetDataEnd(ident_t *loc, DeviceTy &Device, int32_t ArgNum,
bool ForceDelete = ArgTypes[I] & OMP_TGT_MAPTYPE_DELETE;
bool HasCloseModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_CLOSE;
bool HasPresentModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_PRESENT;
bool HasHoldModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_OMPX_HOLD;
// If PTR_AND_OBJ, HstPtrBegin is address of pointee
void *TgtPtrBegin =
Device.getTgtPtrBegin(HstPtrBegin, DataSize, IsLast, UpdateRef,
IsHostPtr, !IsImplicit, ForceDelete);
void *TgtPtrBegin = Device.getTgtPtrBegin(
HstPtrBegin, DataSize, IsLast, UpdateRef, HasHoldModifier, IsHostPtr,
!IsImplicit, ForceDelete);
if (!TgtPtrBegin && (DataSize || HasPresentModifier)) {
DP("Mapping does not exist (%s)\n",
(HasPresentModifier ? "'present' map type modifier" : "ignored"));
@@ -799,7 +806,8 @@ int targetDataEnd(ident_t *loc, DeviceTy &Device, int32_t ArgNum,
// Add pointer to the buffer for later deallocation
if (DelEntry)
DeallocTgtPtrs.emplace_back(HstPtrBegin, DataSize, HasCloseModifier);
DeallocTgtPtrs.emplace_back(HstPtrBegin, DataSize, HasCloseModifier,
HasHoldModifier);
}
}
@@ -816,7 +824,7 @@ int targetDataEnd(ident_t *loc, DeviceTy &Device, int32_t ArgNum,
if (FromMapperBase && FromMapperBase == Info.HstPtrBegin)
continue;
Ret = Device.deallocTgtPtr(Info.HstPtrBegin, Info.DataSize,
Info.HasCloseModifier);
Info.HasCloseModifier, Info.HasHoldModifier);
if (Ret != OFFLOAD_SUCCESS) {
REPORT("Deallocating data from device failed.\n");
return OFFLOAD_FAIL;
@@ -831,8 +839,9 @@ static int targetDataContiguous(ident_t *loc, DeviceTy &Device, void *ArgsBase,
int64_t ArgType, AsyncInfoTy &AsyncInfo) {
TIMESCOPE_WITH_IDENT(loc);
bool IsLast, IsHostPtr;
void *TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBegin, ArgSize, IsLast, false,
IsHostPtr, /*MustContain=*/true);
void *TgtPtrBegin = Device.getTgtPtrBegin(
HstPtrBegin, ArgSize, IsLast, /*UpdateRefCount=*/false,
/*UseHoldRefCount=*/false, IsHostPtr, /*MustContain=*/true);
if (!TgtPtrBegin) {
DP("hst data:" DPxMOD " not found, becomes a noop\n", DPxPTR(HstPtrBegin));
if (ArgType & OMP_TGT_MAPTYPE_PRESENT) {
@@ -1291,8 +1300,9 @@ static int processDataBefore(ident_t *loc, int64_t DeviceId, void *HostPtr,
uint64_t Delta = (uint64_t)HstPtrBegin - (uint64_t)HstPtrBase;
void *TgtPtrBegin = (void *)((uintptr_t)TgtPtrBase + Delta);
void *&PointerTgtPtrBegin = AsyncInfo.getVoidPtrLocation();
PointerTgtPtrBegin = Device.getTgtPtrBegin(HstPtrVal, ArgSizes[I],
IsLast, false, IsHostPtr);
PointerTgtPtrBegin = Device.getTgtPtrBegin(
HstPtrVal, ArgSizes[I], IsLast, /*UpdateRefCount=*/false,
/*UseHoldRefCount=*/false, IsHostPtr);
if (!PointerTgtPtrBegin) {
DP("No lambda captured variable mapped (" DPxMOD ") - ignored\n",
DPxPTR(HstPtrVal));
@@ -1348,7 +1358,8 @@ static int processDataBefore(ident_t *loc, int64_t DeviceId, void *HostPtr,
if (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ)
HstPtrBase = *reinterpret_cast<void **>(HstPtrBase);
TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBegin, ArgSizes[I], IsLast,
false, IsHostPtr);
/*UpdateRefCount=*/false,
/*UseHoldRefCount=*/false, IsHostPtr);
TgtBaseOffset = (intptr_t)HstPtrBase - (intptr_t)HstPtrBegin;
#ifdef OMPTARGET_DEBUG
void *TgtPtrBase = (void *)((intptr_t)TgtPtrBegin + TgtBaseOffset);

10
openmp/libomptarget/src/private.h
View File

@@ -111,16 +111,18 @@ static inline void dumpTargetPointerMappings(const ident_t *Loc,
INFO(OMP_INFOTYPE_ALL, Device.DeviceID,
"OpenMP Host-Device pointer mappings after block at %s:%d:%d:\n",
Kernel.getFilename(), Kernel.getLine(), Kernel.getColumn());
INFO(OMP_INFOTYPE_ALL, Device.DeviceID, "%-18s %-18s %s %s %s\n", "Host Ptr",
"Target Ptr", "Size (B)", "RefCount", "Declaration");
INFO(OMP_INFOTYPE_ALL, Device.DeviceID, "%-18s %-18s %s %s %s %s\n",
"Host Ptr", "Target Ptr", "Size (B)", "DynRefCount", "HoldRefCount",
"Declaration");
Device.DataMapMtx.lock();
for (const auto &HostTargetMap : Device.HostDataToTargetMap) {
SourceInfo Info(HostTargetMap.HstPtrName);
INFO(OMP_INFOTYPE_ALL, Device.DeviceID,
DPxMOD " " DPxMOD " %-8" PRIuPTR " %-8s %s at %s:%d:%d\n",
DPxMOD " " DPxMOD " %-8" PRIuPTR " %-11s %-12s %s at %s:%d:%d\n",
DPxPTR(HostTargetMap.HstPtrBegin), DPxPTR(HostTargetMap.TgtPtrBegin),
HostTargetMap.HstPtrEnd - HostTargetMap.HstPtrBegin,
HostTargetMap.refCountToStr().c_str(), Info.getName(),
HostTargetMap.dynRefCountToStr().c_str(),
HostTargetMap.holdRefCountToStr().c_str(), Info.getName(),
Info.getFilename(), Info.getLine(), Info.getColumn());
}
Device.DataMapMtx.unlock();

68
openmp/libomptarget/test/mapping/ompx_hold/omp_target_disassociate_ptr.c Normal file
View File

@@ -0,0 +1,68 @@
// omp_target_disassociate_ptr should always fail if the hold reference count is
// non-zero, regardless of the dynamic reference count. When the latter is
// finite, the implementation happens to choose to report the hold diagnostic.
// RUN: %libomptarget-compile-generic -fopenmp-extensions
// RUN: %not %libomptarget-run-generic 0 2>&1 | %fcheck-generic
// RUN: %not %libomptarget-run-generic 1 2>&1 | %fcheck-generic
// RUN: %not %libomptarget-run-generic inf 2>&1 | %fcheck-generic
// RUN: %libomptarget-compile-generic -fopenmp-extensions -DHOLD_MORE
// RUN: %not %libomptarget-run-generic 0 2>&1 | %fcheck-generic
// RUN: %not %libomptarget-run-generic 1 2>&1 | %fcheck-generic
// RUN: %not %libomptarget-run-generic inf 2>&1 | %fcheck-generic
#include <omp.h>
#include <stdio.h>
#include <limits.h>
#include <string.h>
int main(int argc, char *argv[]) {
// Parse command line.
int DynRef;
if (argc != 2) {
fprintf(stderr, "bad arguments\n");
return 1;
}
if (0 == strcmp(argv[1], "inf"))
DynRef = INT_MAX;
else
DynRef = atoi(argv[1]);
// Allocate and set dynamic reference count as specified.
int DevNum = omp_get_default_device();
int X;
void *XDev = omp_target_alloc(sizeof X, DevNum);
if (!XDev) {
fprintf(stderr, "omp_target_alloc failed\n");
return 1;
}
if (DynRef == INT_MAX) {
if (omp_target_associate_ptr(&X, &XDev, sizeof X, 0, DevNum)) {
fprintf(stderr, "omp_target_associate_ptr failed\n");
return 1;
}
} else {
for (int I = 0; I < DynRef; ++I) {
#pragma omp target enter data map(alloc: X)
}
}
// Disassociate while hold reference count > 0.
int Status = 0;
#pragma omp target data map(ompx_hold,alloc: X)
#if HOLD_MORE
#pragma omp target data map(ompx_hold,alloc: X)
#pragma omp target data map(ompx_hold,alloc: X)
#endif
{
// CHECK: Libomptarget error: Trying to disassociate a pointer with a
// CHECK-SAME: non-zero hold reference count
// CHECK-NEXT: omp_target_disassociate_ptr failed
if (omp_target_disassociate_ptr(&X, DevNum)) {
fprintf(stderr, "omp_target_disassociate_ptr failed\n");
Status = 1;
}
}
return Status;
}

202
openmp/libomptarget/test/mapping/ompx_hold/struct.c Normal file
View File

@@ -0,0 +1,202 @@
// RUN: %libomptarget-compile-generic -fopenmp-extensions
// RUN: %libomptarget-run-generic | %fcheck-generic -strict-whitespace
#include <omp.h>
#include <stdio.h>
#define CHECK_PRESENCE(Var1, Var2, Var3) \
printf(" presence of %s, %s, %s: %d, %d, %d\n", \
#Var1, #Var2, #Var3, \
omp_target_is_present(&(Var1), omp_get_default_device()), \
omp_target_is_present(&(Var2), omp_get_default_device()), \
omp_target_is_present(&(Var3), omp_get_default_device()))
#define CHECK_VALUES(Var1, Var2) \
printf(" values of %s, %s: %d, %d\n", \
#Var1, #Var2, (Var1), (Var2))
int main() {
struct S { int i; int j; } s;
// CHECK: presence of s, s.i, s.j: 0, 0, 0
CHECK_PRESENCE(s, s.i, s.j);
// =======================================================================
// Check that ompx_hold keeps entire struct present.
// -----------------------------------------------------------------------
// CHECK-LABEL: check:{{.*}}
printf("check: ompx_hold only on first member\n");
s.i = 20;
s.j = 30;
#pragma omp target data map(tofrom: s) map(ompx_hold,tofrom: s.i) \
map(tofrom: s.j)
{
// CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
CHECK_PRESENCE(s, s.i, s.j);
#pragma omp target map(tofrom: s)
{
s.i = 21;
s.j = 31;
}
#pragma omp target exit data map(delete: s, s.i)
// ompx_hold on s.i applies to all of s.
// CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
// CHECK-NEXT: values of s.i, s.j: 20, 30
CHECK_PRESENCE(s, s.i, s.j);
CHECK_VALUES(s.i, s.j);
}
// CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0
// CHECK-NEXT: values of s.i, s.j: 21, 31
CHECK_PRESENCE(s, s.i, s.j);
CHECK_VALUES(s.i, s.j);
// -----------------------------------------------------------------------
// CHECK-LABEL: check:{{.*}}
printf("check: ompx_hold only on last member\n");
s.i = 20;
s.j = 30;
#pragma omp target data map(tofrom: s) map(tofrom: s.i) \
map(ompx_hold,tofrom: s.j)
{
// CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
CHECK_PRESENCE(s, s.i, s.j);
#pragma omp target map(tofrom: s)
{
s.i = 21;
s.j = 31;
}
#pragma omp target exit data map(delete: s, s.i)
// ompx_hold on s.j applies to all of s.
// CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
// CHECK-NEXT: values of s.i, s.j: 20, 30
CHECK_PRESENCE(s, s.i, s.j);
CHECK_VALUES(s.i, s.j);
}
// CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0
// CHECK-NEXT: values of s.i, s.j: 21, 31
CHECK_PRESENCE(s, s.i, s.j);
CHECK_VALUES(s.i, s.j);
// -----------------------------------------------------------------------
// CHECK-LABEL: check:{{.*}}
printf("check: ompx_hold only on struct\n");
s.i = 20;
s.j = 30;
#pragma omp target data map(ompx_hold,tofrom: s) map(tofrom: s.i) \
map(tofrom: s.j)
{
// CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
CHECK_PRESENCE(s, s.i, s.j);
#pragma omp target map(tofrom: s)
{
s.i = 21;
s.j = 31;
}
#pragma omp target exit data map(delete: s, s.i)
// CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
// CHECK-NEXT: values of s.i, s.j: 20, 30
CHECK_PRESENCE(s, s.i, s.j);
CHECK_VALUES(s.i, s.j);
}
// CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0
// CHECK-NEXT: values of s.i, s.j: 21, 31
CHECK_PRESENCE(s, s.i, s.j);
CHECK_VALUES(s.i, s.j);
// =======================================================================
// Check that transfer to/from host checks reference count correctly.
// -----------------------------------------------------------------------
// CHECK-LABEL: check:{{.*}}
printf("check: parent DynRefCount=1 is not sufficient for transfer\n");
s.i = 20;
s.j = 30;
#pragma omp target data map(ompx_hold, tofrom: s)
#pragma omp target data map(ompx_hold, tofrom: s)
{
// CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
CHECK_PRESENCE(s, s.i, s.j);
#pragma omp target map(from: s.i, s.j)
{
s.i = 21;
s.j = 31;
} // No transfer here even though parent's DynRefCount=1.
// CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
// CHECK-NEXT: values of s.i, s.j: 20, 30
CHECK_PRESENCE(s, s.i, s.j);
CHECK_VALUES(s.i, s.j);
#pragma omp target map(to: s.i, s.j)
{ // No transfer here even though parent's DynRefCount=1.
// CHECK-NEXT: values of s.i, s.j: 21, 31
CHECK_VALUES(s.i, s.j);
}
}
// CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0
// CHECK-NEXT: values of s.i, s.j: 21, 31
CHECK_PRESENCE(s, s.i, s.j);
CHECK_VALUES(s.i, s.j);
// -----------------------------------------------------------------------
// CHECK-LABEL: check:{{.*}}
printf("check: parent HoldRefCount=1 is not sufficient for transfer\n");
s.i = 20;
s.j = 30;
#pragma omp target data map(tofrom: s)
#pragma omp target data map(tofrom: s)
{
// CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
CHECK_PRESENCE(s, s.i, s.j);
#pragma omp target map(ompx_hold, from: s.i, s.j)
{
s.i = 21;
s.j = 31;
} // No transfer here even though parent's HoldRefCount=1.
// CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
// CHECK-NEXT: values of s.i, s.j: 20, 30
CHECK_PRESENCE(s, s.i, s.j);
CHECK_VALUES(s.i, s.j);
#pragma omp target map(ompx_hold, to: s.i, s.j)
{ // No transfer here even though parent's HoldRefCount=1.
// CHECK-NEXT: values of s.i, s.j: 21, 31
CHECK_VALUES(s.i, s.j);
}
}
// CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0
// CHECK-NEXT: values of s.i, s.j: 21, 31
CHECK_PRESENCE(s, s.i, s.j);
CHECK_VALUES(s.i, s.j);
// -----------------------------------------------------------------------
// CHECK-LABEL: check:{{.*}}
//
// At the beginning of a region, if the parent's TotalRefCount=1, then the
// transfer should happen.
//
// At the end of a region, it also must be true that the reference count being
// decremented is the reference count that is 1.
printf("check: parent TotalRefCount=1 is not sufficient for transfer\n");
s.i = 20;
s.j = 30;
#pragma omp target data map(ompx_hold, tofrom: s)
{
// CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
CHECK_PRESENCE(s, s.i, s.j);
#pragma omp target map(ompx_hold, tofrom: s.i, s.j)
{
s.i = 21;
s.j = 31;
}
#pragma omp target exit data map(from: s.i, s.j)
// No transfer here even though parent's TotalRefCount=1.
// CHECK-NEXT: presence of s, s.i, s.j: 1, 1, 1
// CHECK-NEXT: values of s.i, s.j: 20, 30
CHECK_PRESENCE(s, s.i, s.j);
CHECK_VALUES(s.i, s.j);
}
// CHECK-NEXT: presence of s, s.i, s.j: 0, 0, 0
// CHECK-NEXT: values of s.i, s.j: 21, 31
CHECK_PRESENCE(s, s.i, s.j);
CHECK_VALUES(s.i, s.j);
return 0;
}

236
openmp/libomptarget/test/mapping/ompx_hold/target-data.c Normal file
View File

@@ -0,0 +1,236 @@
// RUN: %libomptarget-compile-generic -fopenmp-extensions
// RUN: %libomptarget-run-generic | %fcheck-generic -strict-whitespace
#include <omp.h>
#include <stdio.h>
#define CHECK_PRESENCE(Var1, Var2, Var3) \
printf(" presence of %s, %s, %s: %d, %d, %d\n", \
#Var1, #Var2, #Var3, \
omp_target_is_present(&Var1, omp_get_default_device()), \
omp_target_is_present(&Var2, omp_get_default_device()), \
omp_target_is_present(&Var3, omp_get_default_device()))
int main() {
int m, r, d;
// CHECK: presence of m, r, d: 0, 0, 0
CHECK_PRESENCE(m, r, d);
// -----------------------------------------------------------------------
// CHECK-NEXT: check:{{.*}}
printf("check: dyn>0, hold=0, dec/reset dyn=0\n");
// CHECK-NEXT: structured{{.*}}
printf(" structured dec of dyn\n");
#pragma omp target data map(tofrom: m) map(alloc: r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
#pragma omp target data map(tofrom: m) map(alloc: r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
}
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
}
// CHECK-NEXT: presence of m, r, d: 0, 0, 0
CHECK_PRESENCE(m, r, d);
// CHECK-NEXT: dynamic{{.*}}
printf(" dynamic dec/reset of dyn\n");
#pragma omp target data map(tofrom: m) map(alloc: r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
#pragma omp target data map(tofrom: m) map(alloc: r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
#pragma omp target exit data map(from: m) map(release: r)
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
#pragma omp target exit data map(from: m) map(release: r) map(delete: d)
// CHECK-NEXT: presence of m, r, d: 0, 0, 0
CHECK_PRESENCE(m, r, d);
}
// CHECK-NEXT: presence of m, r, d: 0, 0, 0
CHECK_PRESENCE(m, r, d);
#pragma omp target exit data map(from: m) map(release: r) map(delete: d)
// CHECK-NEXT: presence of m, r, d: 0, 0, 0
CHECK_PRESENCE(m, r, d);
}
// CHECK-NEXT: presence of m, r, d: 0, 0, 0
CHECK_PRESENCE(m, r, d);
// -----------------------------------------------------------------------
// CHECK: check:{{.*}}
printf("check: dyn=0, hold>0, dec/reset dyn=0, dec hold=0\n");
// Structured dec of dyn would require dyn>0.
// CHECK-NEXT: dynamic{{.*}}
printf(" dynamic dec/reset of dyn\n");
#pragma omp target data map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
#pragma omp target data map(ompx_hold, tofrom: m) \
map(ompx_hold, alloc: r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
#pragma omp target exit data map(from: m) map(release: r)
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
#pragma omp target exit data map(from: m) map(release: r) map(delete: d)
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
}
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
#pragma omp target exit data map(from: m) map(release: r) map(delete: d)
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
}
// CHECK-NEXT: presence of m, r, d: 0, 0, 0
CHECK_PRESENCE(m, r, d);
// -----------------------------------------------------------------------
// CHECK: check:{{.*}}
printf("check: dyn>0, hold>0, dec/reset dyn=0, dec hold=0\n");
// CHECK-NEXT: structured{{.*}}
printf(" structured dec of dyn\n");
#pragma omp target data map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
#pragma omp target data map(ompx_hold, tofrom: m) \
map(ompx_hold, alloc: r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
#pragma omp target data map(tofrom: m) map(alloc: r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
#pragma omp target data map(tofrom: m) map(alloc: r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
}
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
}
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
}
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
}
// CHECK-NEXT: presence of m, r, d: 0, 0, 0
CHECK_PRESENCE(m, r, d);
// CHECK-NEXT: dynamic{{.*}}
printf(" dynamic dec/reset of dyn\n");
#pragma omp target enter data map(to: m) map(alloc: r, d)
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
#pragma omp target enter data map(to: m) map(alloc: r, d)
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
#pragma omp target data map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
#pragma omp target data map(ompx_hold, tofrom: m) \
map(ompx_hold, alloc: r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
#pragma omp target exit data map(from: m) map(release: r)
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
#pragma omp target exit data map(from: m) map(release: r) map(delete: d)
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
}
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
#pragma omp target exit data map(from: m) map(release: r) map(delete: d)
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
}
// CHECK-NEXT: presence of m, r, d: 0, 0, 0
CHECK_PRESENCE(m, r, d);
// -----------------------------------------------------------------------
// CHECK: check:{{.*}}
printf("check: dyn>0, hold>0, dec hold=0, dec/reset dyn=0\n");
// CHECK-NEXT: structured{{.*}}
printf(" structured dec of dyn\n");
#pragma omp target data map(tofrom: m) map(alloc: r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
#pragma omp target data map(tofrom: m) map(alloc: r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
#pragma omp target data map(ompx_hold, tofrom: m) \
map(ompx_hold, alloc: r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
#pragma omp target data map(ompx_hold, tofrom: m) \
map(ompx_hold, alloc: r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
}
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
}
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
}
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
}
// CHECK-NEXT: presence of m, r, d: 0, 0, 0
CHECK_PRESENCE(m, r, d);
// CHECK-NEXT: dynamic{{.*}}
printf(" dynamic dec/reset of dyn\n");
#pragma omp target enter data map(to: m) map(alloc: r, d)
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
#pragma omp target enter data map(to: m) map(alloc: r, d)
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
#pragma omp target data map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
#pragma omp target data map(ompx_hold, tofrom: m) \
map(ompx_hold, alloc: r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
}
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
}
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
#pragma omp target exit data map(from: m) map(release: r)
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
#pragma omp target exit data map(from: m) map(release: r) map(delete: d)
// CHECK-NEXT: presence of m, r, d: 0, 0, 0
CHECK_PRESENCE(m, r, d);
return 0;
}

164
openmp/libomptarget/test/mapping/ompx_hold/target.c Normal file
View File

@@ -0,0 +1,164 @@
// RUN: %libomptarget-compile-generic -fopenmp-extensions
// RUN: %libomptarget-run-generic | %fcheck-generic -strict-whitespace
#include <omp.h>
#include <stdio.h>
#define CHECK_PRESENCE(Var1, Var2, Var3) \
printf(" presence of %s, %s, %s: %d, %d, %d\n", \
#Var1, #Var2, #Var3, \
omp_target_is_present(&Var1, omp_get_default_device()), \
omp_target_is_present(&Var2, omp_get_default_device()), \
omp_target_is_present(&Var3, omp_get_default_device()))
int main() {
int m, r, d;
// CHECK: presence of m, r, d: 0, 0, 0
CHECK_PRESENCE(m, r, d);
// -----------------------------------------------------------------------
// CHECK-NEXT: check:{{.*}}
printf("check: dyn>0, hold=0, dec dyn=0\n");
// CHECK-NEXT: once
printf(" once\n");
#pragma omp target map(tofrom: m) map(alloc: r, d)
;
// CHECK-NEXT: presence of m, r, d: 0, 0, 0
CHECK_PRESENCE(m, r, d);
// CHECK-NEXT: twice
printf(" twice\n");
#pragma omp target data map(tofrom: m) map(alloc: r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
#pragma omp target map(tofrom: m) map(alloc: r, d)
;
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
}
// CHECK-NEXT: presence of m, r, d: 0, 0, 0
CHECK_PRESENCE(m, r, d);
// -----------------------------------------------------------------------
// CHECK: check:{{.*}}
printf("check: dyn=0, hold>0, dec hold=0\n");
// CHECK-NEXT: once
printf(" once\n");
#pragma omp target map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
;
// CHECK-NEXT: presence of m, r, d: 0, 0, 0
CHECK_PRESENCE(m, r, d);
// CHECK-NEXT: twice
printf(" twice\n");
#pragma omp target data map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
#pragma omp target map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
;
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
}
// CHECK-NEXT: presence of m, r, d: 0, 0, 0
CHECK_PRESENCE(m, r, d);
// -----------------------------------------------------------------------
// CHECK: check:{{.*}}
printf("check: dyn>0, hold>0, dec dyn=0, dec hold=0\n");
// CHECK-NEXT: once each
printf(" once each\n");
#pragma omp target data map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
#pragma omp target map(tofrom: m) map(alloc: r, d)
;
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
}
// CHECK-NEXT: presence of m, r, d: 0, 0, 0
CHECK_PRESENCE(m, r, d);
// CHECK-NEXT: twice each
printf(" twice each\n");
#pragma omp target data map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
#pragma omp target data map(ompx_hold, tofrom: m) \
map(ompx_hold, alloc: r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
#pragma omp target data map(tofrom: m) map(alloc: r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
#pragma omp target map(tofrom: m) map(alloc: r, d)
;
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
}
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
}
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
}
// CHECK-NEXT: presence of m, r, d: 0, 0, 0
CHECK_PRESENCE(m, r, d);
// -----------------------------------------------------------------------
// CHECK: check:{{.*}}
printf("check: dyn>0, hold>0, dec hold=0, dec dyn=0\n");
// CHECK-NEXT: once each
printf(" once each\n");
#pragma omp target data map(tofrom: m) map(alloc: r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
#pragma omp target map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
;
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
}
// CHECK-NEXT: presence of m, r, d: 0, 0, 0
CHECK_PRESENCE(m, r, d);
// CHECK-NEXT: twice each
printf(" twice each\n");
#pragma omp target data map(tofrom: m) map(alloc: r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
#pragma omp target data map(tofrom: m) map(alloc: r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
#pragma omp target data map(ompx_hold, tofrom: m) \
map(ompx_hold, alloc: r, d)
{
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
#pragma omp target map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
;
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
}
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
}
// CHECK-NEXT: presence of m, r, d: 1, 1, 1
CHECK_PRESENCE(m, r, d);
}
// CHECK-NEXT: presence of m, r, d: 0, 0, 0
CHECK_PRESENCE(m, r, d);
return 0;
}

33
openmp/libomptarget/test/offloading/info.c
View File

@@ -1,4 +1,7 @@
// RUN: %libomptarget-compile-nvptx64-nvidia-cuda -gline-tables-only && env LIBOMPTARGET_INFO=63 %libomptarget-run-nvptx64-nvidia-cuda 2>&1 | %fcheck-nvptx64-nvidia-cuda -allow-empty -check-prefix=INFO
// RUN: %libomptarget-compile-nvptx64-nvidia-cuda \
// RUN: -gline-tables-only -fopenmp-extensions
// RUN: env LIBOMPTARGET_INFO=63 %libomptarget-run-nvptx64-nvidia-cuda 2>&1 | \
// RUN: %fcheck-nvptx64-nvidia-cuda -allow-empty -check-prefix=INFO
// REQUIRES: nvptx64-nvidia-cuda
#include <stdio.h>
@@ -23,24 +26,24 @@ int main() {
// INFO: Libomptarget device 0 info: alloc(A[0:64])[256]
// INFO: Libomptarget device 0 info: tofrom(B[0:64])[256]
// INFO: Libomptarget device 0 info: to(C[0:64])[256]
// INFO: Libomptarget device 0 info: Creating new map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, RefCount=1, Name=A[0:64]
// INFO: Libomptarget device 0 info: Creating new map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, RefCount=1, Name=B[0:64]
// INFO: Libomptarget device 0 info: Creating new map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, DynRefCount=1, HoldRefCount=0, Name=A[0:64]
// INFO: Libomptarget device 0 info: Creating new map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, DynRefCount=0, HoldRefCount=1, Name=B[0:64]
// INFO: Libomptarget device 0 info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=256, Name=B[0:64]
// INFO: Libomptarget device 0 info: Creating new map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, RefCount=1, Name=C[0:64]
// INFO: Libomptarget device 0 info: Creating new map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, DynRefCount=1, HoldRefCount=0, Name=C[0:64]
// INFO: Libomptarget device 0 info: Copying data from host to device, HstPtr={{.*}}, TgtPtr={{.*}}, Size=256, Name=C[0:64]
// INFO: Libomptarget device 0 info: OpenMP Host-Device pointer mappings after block at info.c:{{[0-9]+}}:{{[0-9]+}}:
// INFO: Libomptarget device 0 info: Host Ptr Target Ptr Size (B) RefCount Declaration
// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256 1 C[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256 1 B[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256 1 A[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
// INFO: Libomptarget device 0 info: Host Ptr Target Ptr Size (B) DynRefCount HoldRefCount Declaration
// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256 1 0 C[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256 0 1 B[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256 1 0 A[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
// INFO: Libomptarget device 0 info: Entering OpenMP kernel at info.c:{{[0-9]+}}:{{[0-9]+}} with 1 arguments:
// INFO: Libomptarget device 0 info: firstprivate(val)[4]
// INFO: CUDA device 0 info: Launching kernel __omp_offloading_{{.*}}main{{.*}} with {{[0-9]+}} blocks and {{[0-9]+}} threads in Generic mode
// INFO: Libomptarget device 0 info: OpenMP Host-Device pointer mappings after block at info.c:{{[0-9]+}}:{{[0-9]+}}:
// INFO: Libomptarget device 0 info: Host Ptr Target Ptr Size (B) RefCount Declaration
// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256 1 C[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256 1 B[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256 1 A[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
// INFO: Libomptarget device 0 info: Host Ptr Target Ptr Size (B) DynRefCount HoldRefCount Declaration
// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256 1 0 C[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256 0 1 B[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
// INFO: Libomptarget device 0 info: {{.*}} {{.*}} 256 1 0 A[0:64] at info.c:{{[0-9]+}}:{{[0-9]+}}
// INFO: Libomptarget device 0 info: Exiting OpenMP data region at info.c:{{[0-9]+}}:{{[0-9]+}} with 3 arguments:
// INFO: Libomptarget device 0 info: alloc(A[0:64])[256]
// INFO: Libomptarget device 0 info: tofrom(B[0:64])[256]
@@ -50,9 +53,9 @@ int main() {
// INFO: Libomptarget device 0 info: Removing map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, Name=B[0:64]
// INFO: Libomptarget device 0 info: Removing map entry with HstPtrBegin={{.*}}, TgtPtrBegin={{.*}}, Size=256, Name=A[0:64]
// INFO: Libomptarget device 0 info: OpenMP Host-Device pointer mappings after block at info.c:[[#%u,]]:[[#%u,]]:
// INFO: Libomptarget device 0 info: Host Ptr Target Ptr Size (B) RefCount Declaration
// INFO: Libomptarget device 0 info: [[#%#x,]] [[#%#x,]] 4 INF unknown at unknown:0:0
#pragma omp target data map(alloc:A[0:N]) map(tofrom:B[0:N]) map(to:C[0:N])
// INFO: Libomptarget device 0 info: Host Ptr Target Ptr Size (B) DynRefCount HoldRefCount Declaration
// INFO: Libomptarget device 0 info: [[#%#x,]] [[#%#x,]] 4 INF 0 unknown at unknown:0:0
#pragma omp target data map(alloc:A[0:N]) map(ompx_hold,tofrom:B[0:N]) map(to:C[0:N])
#pragma omp target firstprivate(val)
{ val = 1; }