| //===-- OpenMP/Mapping.cpp - OpenMP/OpenACC pointer mapping impl. ---------===// | |
| // | |
| // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. | |
| // See https://llvm.org/LICENSE.txt for license information. | |
| // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception | |
| // | |
| //===----------------------------------------------------------------------===// | |
| // | |
| //===----------------------------------------------------------------------===// | |
| #include "OpenMP/Mapping.h" | |
| #include "PluginManager.h" | |
| #include "Shared/Debug.h" | |
| #include "Shared/Requirements.h" | |
| #include "device.h" | |
| using namespace llvm::omp::target::debug; | |
| /// Dump a table of all the host-target pointer pairs on failure | |
| void dumpTargetPointerMappings(const ident_t *Loc, DeviceTy &Device, | |
| bool toStdOut) { | |
| MappingInfoTy::HDTTMapAccessorTy HDTTMap = | |
| Device.getMappingInfo().HostDataToTargetMap.getExclusiveAccessor(); | |
| if (HDTTMap->empty()) { | |
| DUMP_INFO(toStdOut, OMP_INFOTYPE_ALL, Device.DeviceID, | |
| "OpenMP Host-Device pointer mappings table empty\n"); | |
| return; | |
| } | |
| SourceInfo Kernel(Loc); | |
| DUMP_INFO(toStdOut, OMP_INFOTYPE_ALL, Device.DeviceID, | |
| "OpenMP Host-Device pointer mappings after block at %s:%d:%d:\n", | |
| Kernel.getFilename(), Kernel.getLine(), Kernel.getColumn()); | |
| DUMP_INFO(toStdOut, OMP_INFOTYPE_ALL, Device.DeviceID, | |
| "%-18s %-18s %s %s %s %s\n", "Host Ptr", "Target Ptr", "Size (B)", | |
| "DynRefCount", "HoldRefCount", "Declaration"); | |
| for (const auto &It : *HDTTMap) { | |
| HostDataToTargetTy &HDTT = *It.HDTT; | |
| SourceInfo Info(HDTT.HstPtrName); | |
| DUMP_INFO(toStdOut, OMP_INFOTYPE_ALL, Device.DeviceID, | |
| DPxMOD " " DPxMOD " %-8" PRIuPTR " %-11s %-12s %s at %s:%d:%d\n", | |
| DPxPTR(HDTT.HstPtrBegin), DPxPTR(HDTT.TgtPtrBegin), | |
| HDTT.HstPtrEnd - HDTT.HstPtrBegin, | |
| HDTT.dynRefCountToStr().c_str(), HDTT.holdRefCountToStr().c_str(), | |
| Info.getName(), Info.getFilename(), Info.getLine(), | |
| Info.getColumn()); | |
| } | |
| } | |
| int MappingInfoTy::associatePtr(void *HstPtrBegin, void *TgtPtrBegin, | |
| int64_t Size) { | |
| HDTTMapAccessorTy HDTTMap = HostDataToTargetMap.getExclusiveAccessor(); | |
| // Check if entry exists | |
| auto It = HDTTMap->find(HstPtrBegin); | |
| if (It != HDTTMap->end()) { | |
| HostDataToTargetTy &HDTT = *It->HDTT; | |
| std::lock_guard<HostDataToTargetTy> LG(HDTT); | |
| // Mapping already exists | |
| bool IsValid = HDTT.HstPtrEnd == (uintptr_t)HstPtrBegin + Size && | |
| HDTT.TgtPtrBegin == (uintptr_t)TgtPtrBegin; | |
| if (IsValid) { | |
| ODBG(ODT_MappingExists) << "Attempt to re-associate the same device " | |
| << "ptr+offset with the same " | |
| << "host ptr, nothing to do"; | |
| return OFFLOAD_SUCCESS; | |
| } | |
| REPORT() << "Not allowed to re-associate a different device ptr+offset " | |
| << "with the same host ptr"; | |
| return OFFLOAD_FAIL; | |
| } | |
| // Mapping does not exist, allocate it with refCount=INF | |
| const HostDataToTargetTy &NewEntry = | |
| *HDTTMap | |
| ->emplace(new HostDataToTargetTy( | |
| /*HstPtrBase=*/(uintptr_t)HstPtrBegin, | |
| /*HstPtrBegin=*/(uintptr_t)HstPtrBegin, | |
| /*HstPtrEnd=*/(uintptr_t)HstPtrBegin + Size, | |
| /*TgtAllocBegin=*/(uintptr_t)TgtPtrBegin, | |
| /*TgtPtrBegin=*/(uintptr_t)TgtPtrBegin, | |
| /*UseHoldRefCount=*/false, /*Name=*/nullptr, | |
| /*IsRefCountINF=*/true)) | |
| .first->HDTT; | |
| ODBG(ODT_Mapping) << "Creating new map entry: HstBase=" | |
| << reinterpret_cast<void *>(NewEntry.HstPtrBase) | |
| << ", HstBegin=" | |
| << reinterpret_cast<void *>(NewEntry.HstPtrBegin) | |
| << ", HstEnd=" | |
| << reinterpret_cast<void *>(NewEntry.HstPtrEnd) | |
| << ", TgtBegin=" | |
| << reinterpret_cast<void *>(NewEntry.TgtPtrBegin) | |
| << ", DynRefCount=" << NewEntry.dynRefCountToStr() | |
| << ", HoldRefCount=" << NewEntry.holdRefCountToStr(); | |
| (void)NewEntry; | |
| // Notify the plugin about the new mapping. | |
| return Device.notifyDataMapped(HstPtrBegin, Size); | |
| } | |
| int MappingInfoTy::disassociatePtr(void *HstPtrBegin) { | |
| HDTTMapAccessorTy HDTTMap = HostDataToTargetMap.getExclusiveAccessor(); | |
| auto It = HDTTMap->find(HstPtrBegin); | |
| if (It == HDTTMap->end()) { | |
| REPORT() << "Association not found"; | |
| return OFFLOAD_FAIL; | |
| } | |
| // Mapping exists | |
| HostDataToTargetTy &HDTT = *It->HDTT; | |
| std::lock_guard<HostDataToTargetTy> LG(HDTT); | |
| if (HDTT.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"; | |
| return OFFLOAD_FAIL; | |
| } | |
| if (HDTT.isDynRefCountInf()) { | |
| ODBG(ODT_Mapping) << "Association found, removing it"; | |
| void *Event = HDTT.getEvent(); | |
| delete &HDTT; | |
| if (Event) | |
| Device.destroyEvent(Event); | |
| HDTTMap->erase(It); | |
| return Device.notifyDataUnmapped(HstPtrBegin); | |
| } | |
| REPORT() << "Trying to disassociate a pointer which was not mapped via " | |
| << "omp_target_associate_ptr"; | |
| return OFFLOAD_FAIL; | |
| } | |
| LookupResult MappingInfoTy::lookupMapping(HDTTMapAccessorTy &HDTTMap, | |
| void *HstPtrBegin, int64_t Size, | |
| HostDataToTargetTy *OwnedTPR) { | |
| uintptr_t HP = (uintptr_t)HstPtrBegin; | |
| LookupResult LR; | |
| ODBG(ODT_Mapping) << "Looking up mapping(HstPtrBegin=" << HstPtrBegin | |
| << ", Size=" << Size << ")..."; | |
| if (HDTTMap->empty()) | |
| return LR; | |
| // HDTTMap is std::set, ordered by HstPtrBegin. | |
| // Upper is the first element whose HstPtrBegin > HP. | |
| auto Upper = HDTTMap->upper_bound(HP); | |
| if (Size == 0) { | |
| // HP satisfies | |
| // std::prev(Upper)->HDTT.HstPtrBegin <= HP < Upper->HDTT.HstPtrBegin | |
| if (Upper != HDTTMap->begin()) { | |
| LR.TPR.setEntry(std::prev(Upper)->HDTT, OwnedTPR); | |
| // We know that HP >= LR.TPR.getEntry()->HstPtrBegin | |
| LR.Flags.IsContained = HP < LR.TPR.getEntry()->HstPtrEnd; | |
| } | |
| if (!LR.Flags.IsContained && Upper != HDTTMap->end()) { | |
| LR.TPR.setEntry(Upper->HDTT, OwnedTPR); | |
| // This is a special case: HP is not really contained in the mapped | |
| // address range, but it's contained in the extended address range, | |
| // which suffices to get the mapping of the base pointer. | |
| // We know that HP < LR.TPR.getEntry()->HstPtrBegin | |
| LR.Flags.IsContained = HP >= LR.TPR.getEntry()->HstPtrBase; | |
| } | |
| } else { | |
| if (Upper != HDTTMap->begin()) { | |
| LR.TPR.setEntry(std::prev(Upper)->HDTT, OwnedTPR); | |
| // We know that HP >= LR.TPR.getEntry()->HstPtrBegin | |
| LR.Flags.IsContained = HP < LR.TPR.getEntry()->HstPtrEnd && | |
| (HP + Size) <= LR.TPR.getEntry()->HstPtrEnd; | |
| // Does it extend beyond the mapped address range? | |
| LR.Flags.ExtendsAfter = HP < LR.TPR.getEntry()->HstPtrEnd && | |
| (HP + Size) > LR.TPR.getEntry()->HstPtrEnd; | |
| } | |
| if (!(LR.Flags.IsContained || LR.Flags.ExtendsAfter) && | |
| Upper != HDTTMap->end()) { | |
| LR.TPR.setEntry(Upper->HDTT, OwnedTPR); | |
| // Does it extend into an already mapped address range? | |
| // We know that HP < LR.TPR.getEntry()->HstPtrBegin | |
| LR.Flags.ExtendsBefore = (HP + Size) > LR.TPR.getEntry()->HstPtrBegin; | |
| // Does it extend beyond the mapped address range? | |
| LR.Flags.ExtendsAfter = HP < LR.TPR.getEntry()->HstPtrEnd && | |
| (HP + Size) > LR.TPR.getEntry()->HstPtrEnd; | |
| } | |
| if (LR.Flags.ExtendsBefore) { | |
| ODBG(ODT_Mapping) << "WARNING: Pointer is not mapped but section extends " | |
| << "into already mapped data"; | |
| } | |
| if (LR.Flags.ExtendsAfter) { | |
| ODBG(ODT_Mapping) << "WARNING: Pointer is already mapped but section " | |
| << "extends beyond mapped region"; | |
| } | |
| } | |
| return LR; | |
| } | |
| TargetPointerResultTy MappingInfoTy::getTargetPointer( | |
| HDTTMapAccessorTy &HDTTMap, void *HstPtrBegin, void *HstPtrBase, | |
| int64_t TgtPadding, int64_t Size, map_var_info_t HstPtrName, bool HasFlagTo, | |
| bool HasFlagAlways, bool IsImplicit, bool UpdateRefCount, | |
| bool HasCloseModifier, bool HasPresentModifier, bool HasHoldModifier, | |
| AsyncInfoTy &AsyncInfo, HostDataToTargetTy *OwnedTPR, bool ReleaseHDTTMap, | |
| StateInfoTy *StateInfo) { | |
| LookupResult LR = lookupMapping(HDTTMap, HstPtrBegin, Size, OwnedTPR); | |
| LR.TPR.Flags.IsPresent = true; | |
| // Release the mapping table lock only after the entry is locked by | |
| // attaching it to TPR. Once TPR is destroyed it will release the lock | |
| // on entry. If it is returned the lock will move to the returned object. | |
| // If LR.Entry is already owned/locked we avoid trying to lock it again. | |
| // Check if the pointer is contained. | |
| // If a variable is mapped to the device manually by the user - which would | |
| // lead to the IsContained flag to be true - then we must ensure that the | |
| // device address is returned even under unified memory conditions. | |
| if (LR.Flags.IsContained || | |
| ((LR.Flags.ExtendsBefore || LR.Flags.ExtendsAfter) && IsImplicit)) { | |
| const char *RefCountAction; | |
| if (UpdateRefCount) { | |
| // After this, reference count >= 1. If the reference count was 0 but the | |
| // entry was still there we can reuse the data on the device and avoid a | |
| // new submission. | |
| LR.TPR.getEntry()->incRefCount(HasHoldModifier); | |
| RefCountAction = " (incremented)"; | |
| } else { | |
| // It might have been allocated with the parent, but it's still new. | |
| LR.TPR.Flags.IsNewEntry = LR.TPR.getEntry()->getTotalRefCount() == 1; | |
| RefCountAction = " (update suppressed)"; | |
| } | |
| const char *DynRefCountAction = HasHoldModifier ? "" : RefCountAction; | |
| const char *HoldRefCountAction = HasHoldModifier ? RefCountAction : ""; | |
| uintptr_t Ptr = LR.TPR.getEntry()->TgtPtrBegin + | |
| ((uintptr_t)HstPtrBegin - LR.TPR.getEntry()->HstPtrBegin); | |
| INFO(OMP_INFOTYPE_MAPPING_EXISTS, Device.DeviceID, | |
| "Mapping exists%s with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD | |
| ", Size=%" PRId64 ", DynRefCount=%s%s, HoldRefCount=%s%s, Name=%s\n", | |
| (IsImplicit ? " (implicit)" : ""), DPxPTR(HstPtrBegin), DPxPTR(Ptr), | |
| Size, LR.TPR.getEntry()->dynRefCountToStr().c_str(), DynRefCountAction, | |
| LR.TPR.getEntry()->holdRefCountToStr().c_str(), HoldRefCountAction, | |
| (HstPtrName) ? getNameFromMapping(HstPtrName).c_str() : "unknown"); | |
| LR.TPR.TargetPointer = (void *)Ptr; | |
| } else if ((LR.Flags.ExtendsBefore || LR.Flags.ExtendsAfter) && !IsImplicit) { | |
| // Explicit extension of mapped data - not allowed. | |
| MESSAGE("explicit extension not allowed: host address specified is " DPxMOD | |
| " (%" PRId64 | |
| " bytes), but device allocation maps to host at " DPxMOD | |
| " (%" PRId64 " bytes)", | |
| DPxPTR(HstPtrBegin), Size, DPxPTR(LR.TPR.getEntry()->HstPtrBegin), | |
| LR.TPR.getEntry()->HstPtrEnd - LR.TPR.getEntry()->HstPtrBegin); | |
| if (HasPresentModifier) | |
| MESSAGE("device mapping required by 'present' map type modifier does not " | |
| "exist for host address " DPxMOD " (%" PRId64 " bytes)", | |
| DPxPTR(HstPtrBegin), Size); | |
| } else if ((PM->getRequirements() & OMP_REQ_UNIFIED_SHARED_MEMORY && | |
| !HasCloseModifier) || | |
| (PM->getRequirements() & OMPX_REQ_AUTO_ZERO_COPY)) { | |
| // If unified shared memory is active, implicitly mapped variables that are | |
| // not privatized use host address. Any explicitly mapped variables also use | |
| // host address where correctness is not impeded. In all other cases maps | |
| // are respected. | |
| // In addition to the mapping rules above, the close map modifier forces the | |
| // mapping of the variable to the device. | |
| if (Size) { | |
| INFO(OMP_INFOTYPE_MAPPING_CHANGED, Device.DeviceID, | |
| "Return HstPtrBegin " DPxMOD " Size=%" PRId64 " for unified shared " | |
| "memory\n", | |
| DPxPTR((uintptr_t)HstPtrBegin), Size); | |
| ODBG(ODT_Mapping) << "Return HstPtrBegin " << HstPtrBegin | |
| << " Size=" << Size << " for unified shared memory"; | |
| LR.TPR.Flags.IsPresent = false; | |
| LR.TPR.Flags.IsHostPointer = true; | |
| LR.TPR.TargetPointer = HstPtrBegin; | |
| } | |
| } else if (HasPresentModifier) { | |
| ODBG(ODT_Mapping) << "Mapping required by 'present' map type modifier does " | |
| << "not exist for HstPtrBegin=" << HstPtrBegin | |
| << ", Size=" << Size; | |
| MESSAGE("device mapping required by 'present' map type modifier does not " | |
| "exist for host address " DPxMOD " (%" PRId64 " bytes)", | |
| DPxPTR(HstPtrBegin), Size); | |
| } else if (Size) { | |
| // If it is not contained and Size > 0, we should create a new entry for it. | |
| LR.TPR.Flags.IsNewEntry = true; | |
| uintptr_t TgtAllocBegin = | |
| (uintptr_t)Device.allocData(TgtPadding + Size, HstPtrBegin); | |
| uintptr_t TgtPtrBegin = TgtAllocBegin + TgtPadding; | |
| // Release the mapping table lock only after the entry is locked by | |
| // attaching it to TPR. | |
| LR.TPR.setEntry(HDTTMap | |
| ->emplace(new HostDataToTargetTy( | |
| (uintptr_t)HstPtrBase, (uintptr_t)HstPtrBegin, | |
| (uintptr_t)HstPtrBegin + Size, TgtAllocBegin, | |
| TgtPtrBegin, HasHoldModifier, HstPtrName)) | |
| .first->HDTT); | |
| INFO(OMP_INFOTYPE_MAPPING_CHANGED, Device.DeviceID, | |
| "Creating new map entry with HstPtrBase=" DPxMOD | |
| ", HstPtrBegin=" DPxMOD ", TgtAllocBegin=" DPxMOD | |
| ", TgtPtrBegin=" DPxMOD | |
| ", Size=%ld, DynRefCount=%s, HoldRefCount=%s, Name=%s\n", | |
| DPxPTR(HstPtrBase), DPxPTR(HstPtrBegin), DPxPTR(TgtAllocBegin), | |
| DPxPTR(TgtPtrBegin), Size, | |
| LR.TPR.getEntry()->dynRefCountToStr().c_str(), | |
| LR.TPR.getEntry()->holdRefCountToStr().c_str(), | |
| (HstPtrName) ? getNameFromMapping(HstPtrName).c_str() : "unknown"); | |
| LR.TPR.TargetPointer = (void *)TgtPtrBegin; | |
| // Notify the plugin about the new mapping. | |
| if (Device.notifyDataMapped(HstPtrBegin, Size)) | |
| return TargetPointerResultTy{}; | |
| } else { | |
| // This entry is not present and we did not create a new entry for it. | |
| LR.TPR.Flags.IsPresent = false; | |
| } | |
| // All mapping table modifications have been made. If the user requested it we | |
| // give up the lock. | |
| if (ReleaseHDTTMap) | |
| HDTTMap.destroy(); | |
| // Lambda to check if this pointer was newly allocated on the current region. | |
| // This is needed to handle cases when the TO entry is encountered after an | |
| // alloc entry for the same pointer. In such cases, the ref-count is already | |
| // non-zero when TO is encountered, but we still need to do a transfer. e.g. | |
| // | |
| // struct S { | |
| // int *p; | |
| // }; | |
| // #pragma omp declare mapper(id : S s) map(to: s.p, s.p[0 : 10]) | |
| // | |
| // S s1; | |
| // ... | |
| // #pragma omp target map(alloc : s1.p[0 : 10]) map(mapper(id), to : s1) | |
| auto WasNewlyAllocatedForCurrentRegion = [&]() { | |
| if (!StateInfo) | |
| return false; | |
| bool WasNewlyAllocated = | |
| StateInfo->wasNewlyAllocated(HstPtrBegin).has_value(); | |
| if (WasNewlyAllocated) | |
| ODBG(ODT_Mapping) << "HstPtrBegin " << HstPtrBegin | |
| << " was newly allocated for the current region"; | |
| return WasNewlyAllocated; | |
| }; | |
| // Even if this isn't a new entry, we still need to do a data-transfer if | |
| // the pointer was newly allocated on the current target region. | |
| if (LR.TPR.TargetPointer && !LR.TPR.Flags.IsHostPointer && HasFlagTo && | |
| (LR.TPR.Flags.IsNewEntry || HasFlagAlways || | |
| WasNewlyAllocatedForCurrentRegion()) && | |
| Size != 0) { | |
| // If we have something like: | |
| // #pragma omp target map(to: s.myarr[0:10]) map(to: s.myarr[0:10]) | |
| // then we see two "new" mappings of the struct member s.myarr here -- | |
| // and both have the "IsNewEntry" flag set, so trigger the copy to device | |
| // below. But, the shadow pointer is only initialised on the target for | |
| // the first copy, and the second copy clobbers it. So, this condition | |
| // avoids the (second) copy here if we have already set shadow pointer info. | |
| auto FailOnPtrFound = [HstPtrBegin, Size](ShadowPtrInfoTy &SP) { | |
| if (SP.HstPtrAddr >= HstPtrBegin && | |
| SP.HstPtrAddr < (void *)((char *)HstPtrBegin + Size)) | |
| return OFFLOAD_FAIL; | |
| return OFFLOAD_SUCCESS; | |
| }; | |
| if (LR.TPR.getEntry()->foreachShadowPointerInfo(FailOnPtrFound) == | |
| OFFLOAD_FAIL) { | |
| ODBG(ODT_Mapping) << "Multiple new mappings of " << Size | |
| << " bytes detected (hst:" << HstPtrBegin | |
| << ") -> (tgt:" << LR.TPR.TargetPointer << ")"; | |
| return std::move(LR.TPR); | |
| } | |
| ODBG(ODT_Mapping) << "Moving " << Size << " bytes (hst:" << HstPtrBegin | |
| << ") -> (tgt:" << LR.TPR.TargetPointer << ")"; | |
| int Ret = Device.submitData(LR.TPR.TargetPointer, HstPtrBegin, Size, | |
| AsyncInfo, LR.TPR.getEntry()); | |
| if (Ret != OFFLOAD_SUCCESS) { | |
| REPORT() << "Copying data to device failed."; | |
| // We will also return nullptr if the data movement fails because that | |
| // pointer points to a corrupted memory region so it doesn't make any | |
| // sense to continue to use it. | |
| LR.TPR.TargetPointer = nullptr; | |
| } else if (LR.TPR.getEntry()->addEventIfNecessary(Device, AsyncInfo) != | |
| OFFLOAD_SUCCESS) | |
| return TargetPointerResultTy{}; | |
| } else { | |
| // If not a host pointer and no present modifier, we need to wait for the | |
| // event if it exists. | |
| // Note: Entry might be nullptr because of zero length array section. | |
| if (LR.TPR.getEntry() && !LR.TPR.Flags.IsHostPointer && | |
| !HasPresentModifier) { | |
| void *Event = LR.TPR.getEntry()->getEvent(); | |
| if (Event) { | |
| int Ret = Device.waitEvent(Event, AsyncInfo); | |
| if (Ret != OFFLOAD_SUCCESS) { | |
| // If it fails to wait for the event, we need to return nullptr in | |
| // case of any data race. | |
| REPORT() << "Failed to wait for event " << Event << "."; | |
| return TargetPointerResultTy{}; | |
| } | |
| } | |
| } | |
| } | |
| return std::move(LR.TPR); | |
| } | |
| TargetPointerResultTy MappingInfoTy::getTgtPtrBegin( | |
| void *HstPtrBegin, int64_t Size, bool UpdateRefCount, bool UseHoldRefCount, | |
| bool MustContain, bool ForceDelete, bool FromDataEnd) { | |
| HDTTMapAccessorTy HDTTMap = HostDataToTargetMap.getExclusiveAccessor(); | |
| LookupResult LR = lookupMapping(HDTTMap, HstPtrBegin, Size); | |
| LR.TPR.Flags.IsPresent = true; | |
| if (LR.Flags.IsContained || | |
| (!MustContain && (LR.Flags.ExtendsBefore || LR.Flags.ExtendsAfter))) { | |
| LR.TPR.Flags.IsLast = | |
| LR.TPR.getEntry()->decShouldRemove(UseHoldRefCount, ForceDelete); | |
| if (ForceDelete) { | |
| LR.TPR.getEntry()->resetRefCount(UseHoldRefCount); | |
| assert(LR.TPR.Flags.IsLast == | |
| LR.TPR.getEntry()->decShouldRemove(UseHoldRefCount) && | |
| "expected correct IsLast prediction for reset"); | |
| } | |
| // Increment the number of threads that is using the entry on a | |
| // targetDataEnd, tracking the number of possible "deleters". A thread may | |
| // come to own the entry deletion even if it was not the last one querying | |
| // for it. Thus, we must track every query on targetDataEnds to ensure only | |
| // the last thread that holds a reference to an entry actually deletes it. | |
| if (FromDataEnd) | |
| LR.TPR.getEntry()->incDataEndThreadCount(); | |
| const char *RefCountAction; | |
| if (!UpdateRefCount) { | |
| RefCountAction = " (update suppressed)"; | |
| } else if (LR.TPR.Flags.IsLast) { | |
| LR.TPR.getEntry()->decRefCount(UseHoldRefCount); | |
| assert(LR.TPR.getEntry()->getTotalRefCount() == 0 && | |
| "Expected zero reference count when deletion is scheduled"); | |
| if (ForceDelete) | |
| RefCountAction = " (reset, delayed deletion)"; | |
| else | |
| RefCountAction = " (decremented, delayed deletion)"; | |
| } else { | |
| LR.TPR.getEntry()->decRefCount(UseHoldRefCount); | |
| RefCountAction = " (decremented)"; | |
| } | |
| const char *DynRefCountAction = UseHoldRefCount ? "" : RefCountAction; | |
| const char *HoldRefCountAction = UseHoldRefCount ? RefCountAction : ""; | |
| uintptr_t TP = LR.TPR.getEntry()->TgtPtrBegin + | |
| ((uintptr_t)HstPtrBegin - LR.TPR.getEntry()->HstPtrBegin); | |
| INFO(OMP_INFOTYPE_MAPPING_EXISTS, Device.DeviceID, | |
| "Mapping exists with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD ", " | |
| "Size=%" PRId64 ", DynRefCount=%s%s, HoldRefCount=%s%s\n", | |
| DPxPTR(HstPtrBegin), DPxPTR(TP), Size, | |
| LR.TPR.getEntry()->dynRefCountToStr().c_str(), DynRefCountAction, | |
| LR.TPR.getEntry()->holdRefCountToStr().c_str(), HoldRefCountAction); | |
| LR.TPR.TargetPointer = (void *)TP; | |
| } else if (PM->getRequirements() & OMP_REQ_UNIFIED_SHARED_MEMORY || | |
| PM->getRequirements() & OMPX_REQ_AUTO_ZERO_COPY) { | |
| // If the value isn't found in the mapping and unified shared memory | |
| // is on then it means we have stumbled upon a value which we need to | |
| // use directly from the host. | |
| ODBG(ODT_Mapping) << "Get HstPtrBegin " << HstPtrBegin << " Size=" << Size | |
| << " for unified shared memory"; | |
| LR.TPR.Flags.IsPresent = false; | |
| LR.TPR.Flags.IsHostPointer = true; | |
| LR.TPR.TargetPointer = HstPtrBegin; | |
| } else { | |
| // OpenMP Specification v5.2: if a matching list item is not found, the | |
| // pointer retains its original value as per firstprivate semantics. | |
| LR.TPR.Flags.IsPresent = false; | |
| LR.TPR.Flags.IsHostPointer = false; | |
| LR.TPR.TargetPointer = HstPtrBegin; | |
| } | |
| return std::move(LR.TPR); | |
| } | |
| // Return the target pointer begin (where the data will be moved). | |
| void *MappingInfoTy::getTgtPtrBegin(HDTTMapAccessorTy &HDTTMap, | |
| void *HstPtrBegin, int64_t Size) { | |
| uintptr_t HP = (uintptr_t)HstPtrBegin; | |
| LookupResult LR = lookupMapping(HDTTMap, HstPtrBegin, Size); | |
| if (LR.Flags.IsContained || LR.Flags.ExtendsBefore || LR.Flags.ExtendsAfter) { | |
| uintptr_t TP = | |
| LR.TPR.getEntry()->TgtPtrBegin + (HP - LR.TPR.getEntry()->HstPtrBegin); | |
| return (void *)TP; | |
| } | |
| return NULL; | |
| } | |
| int MappingInfoTy::eraseMapEntry(HDTTMapAccessorTy &HDTTMap, | |
| HostDataToTargetTy *Entry, int64_t Size) { | |
| assert(Entry && "Trying to delete a null entry from the HDTT map."); | |
| assert(Entry->getTotalRefCount() == 0 && | |
| Entry->getDataEndThreadCount() == 0 && | |
| "Trying to delete entry that is in use or owned by another thread."); | |
| INFO(OMP_INFOTYPE_MAPPING_CHANGED, Device.DeviceID, | |
| "Removing map entry with HstPtrBegin=" DPxMOD ", TgtPtrBegin=" DPxMOD | |
| ", Size=%" PRId64 ", Name=%s\n", | |
| DPxPTR(Entry->HstPtrBegin), DPxPTR(Entry->TgtPtrBegin), Size, | |
| (Entry->HstPtrName) ? getNameFromMapping(Entry->HstPtrName).c_str() | |
| : "unknown"); | |
| if (HDTTMap->erase(Entry) == 0) { | |
| REPORT() << "Trying to remove a non-existent map entry"; | |
| return OFFLOAD_FAIL; | |
| } | |
| return OFFLOAD_SUCCESS; | |
| } | |
| int MappingInfoTy::deallocTgtPtrAndEntry(HostDataToTargetTy *Entry, | |
| int64_t Size) { | |
| assert(Entry && "Trying to deallocate a null entry."); | |
| ODBG(ODT_Mapping) << "Deleting tgt data " | |
| << reinterpret_cast<void *>(Entry->TgtPtrBegin) | |
| << " of size " << Size << " by freeing allocation " | |
| << "starting at " | |
| << reinterpret_cast<void *>(Entry->TgtAllocBegin); | |
| void *Event = Entry->getEvent(); | |
| if (Event && Device.destroyEvent(Event) != OFFLOAD_SUCCESS) { | |
| REPORT() << "Failed to destroy event " << Event; | |
| return OFFLOAD_FAIL; | |
| } | |
| int Ret = Device.deleteData((void *)Entry->TgtAllocBegin); | |
| // Notify the plugin about the unmapped memory. | |
| Ret |= Device.notifyDataUnmapped((void *)Entry->HstPtrBegin); | |
| delete Entry; | |
| return Ret; | |
| } | |
| static void printCopyInfoImpl(int DeviceId, bool H2D, void *SrcPtrBegin, | |
| void *DstPtrBegin, int64_t Size, | |
| HostDataToTargetTy *HT) { | |
| INFO(OMP_INFOTYPE_DATA_TRANSFER, DeviceId, | |
| "Copying data from %s to %s, %sPtr=" DPxMOD ", %sPtr=" DPxMOD | |
| ", Size=%" PRId64 ", Name=%s\n", | |
| H2D ? "host" : "device", H2D ? "device" : "host", H2D ? "Hst" : "Tgt", | |
| DPxPTR(H2D ? SrcPtrBegin : DstPtrBegin), H2D ? "Tgt" : "Hst", | |
| DPxPTR(H2D ? DstPtrBegin : SrcPtrBegin), Size, | |
| (HT && HT->HstPtrName) ? getNameFromMapping(HT->HstPtrName).c_str() | |
| : "unknown"); | |
| } | |
| void MappingInfoTy::printCopyInfo( | |
| void *TgtPtrBegin, void *HstPtrBegin, int64_t Size, bool H2D, | |
| HostDataToTargetTy *Entry, MappingInfoTy::HDTTMapAccessorTy *HDTTMapPtr) { | |
| auto HDTTMap = | |
| HostDataToTargetMap.getExclusiveAccessor(!!Entry || !!HDTTMapPtr); | |
| LookupResult LR; | |
| if (!Entry) { | |
| LR = lookupMapping(HDTTMapPtr ? *HDTTMapPtr : HDTTMap, HstPtrBegin, Size); | |
| Entry = LR.TPR.getEntry(); | |
| } | |
| printCopyInfoImpl(Device.DeviceID, H2D, HstPtrBegin, TgtPtrBegin, Size, | |
| Entry); | |
| } |