2196 lines · cpp
1//===------ omptarget.cpp - Target independent OpenMP target RTL -- C++ -*-===//2//3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.4// See https://llvm.org/LICENSE.txt for license information.5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception6//7//===----------------------------------------------------------------------===//8//9// Implementation of the interface to be used by Clang during the codegen of a10// target region.11//12//===----------------------------------------------------------------------===//13 14#include "omptarget.h"15#include "OffloadPolicy.h"16#include "OpenMP/OMPT/Callback.h"17#include "OpenMP/OMPT/Interface.h"18#include "PluginManager.h"19#include "Shared/Debug.h"20#include "Shared/EnvironmentVar.h"21#include "Shared/Utils.h"22#include "device.h"23#include "private.h"24#include "rtl.h"25 26#include "Shared/Profile.h"27 28#include "OpenMP/Mapping.h"29#include "OpenMP/omp.h"30 31#include "llvm/ADT/StringExtras.h"32#include "llvm/ADT/bit.h"33#include "llvm/Frontend/OpenMP/OMPConstants.h"34#include "llvm/Object/ObjectFile.h"35 36#include <cassert>37#include <cstdint>38#include <vector>39 40using llvm::SmallVector;41#ifdef OMPT_SUPPORT42using namespace llvm::omp::target::ompt;43#endif44 45int AsyncInfoTy::synchronize() {46 int Result = OFFLOAD_SUCCESS;47 if (!isQueueEmpty()) {48 switch (SyncType) {49 case SyncTy::BLOCKING:50 // If we have a queue we need to synchronize it now.51 Result = Device.synchronize(*this);52 assert(AsyncInfo.Queue == nullptr &&53 "The device plugin should have nulled the queue to indicate there "54 "are no outstanding actions!");55 break;56 case SyncTy::NON_BLOCKING:57 Result = Device.queryAsync(*this);58 break;59 }60 }61 62 // Run any pending post-processing function registered on this async object.63 if (Result == OFFLOAD_SUCCESS && isQueueEmpty())64 Result = runPostProcessing();65 66 return Result;67}68 69void *&AsyncInfoTy::getVoidPtrLocation() {70 BufferLocations.push_back(nullptr);71 return BufferLocations.back();72}73 74bool AsyncInfoTy::isDone() const { return isQueueEmpty(); }75 76int32_t AsyncInfoTy::runPostProcessing() {77 size_t Size = PostProcessingFunctions.size();78 for (size_t I = 0; I < Size; ++I) {79 const int Result = PostProcessingFunctions[I]();80 if (Result != OFFLOAD_SUCCESS)81 return Result;82 }83 84 // Clear the vector up until the last known function, since post-processing85 // procedures might add new procedures themselves.86 const auto *PrevBegin = PostProcessingFunctions.begin();87 PostProcessingFunctions.erase(PrevBegin, PrevBegin + Size);88 89 return OFFLOAD_SUCCESS;90}91 92bool AsyncInfoTy::isQueueEmpty() const { return AsyncInfo.Queue == nullptr; }93 94/* All begin addresses for partially mapped structs must be aligned, up to 16,95 * in order to ensure proper alignment of members. E.g.96 *97 * struct S {98 * int a; // 4-aligned99 * int b; // 4-aligned100 * int *p; // 8-aligned101 * } s1;102 * ...103 * #pragma omp target map(tofrom: s1.b, s1.p[0:N])104 * {105 * s1.b = 5;106 * for (int i...) s1.p[i] = ...;107 * }108 *109 * Here we are mapping s1 starting from member b, so BaseAddress=&s1=&s1.a and110 * BeginAddress=&s1.b. Let's assume that the struct begins at address 0x100,111 * then &s1.a=0x100, &s1.b=0x104, &s1.p=0x108. Each member obeys the alignment112 * requirements for its type. Now, when we allocate memory on the device, in113 * CUDA's case cuMemAlloc() returns an address which is at least 256-aligned.114 * This means that the chunk of the struct on the device will start at a115 * 256-aligned address, let's say 0x200. Then the address of b will be 0x200 and116 * address of p will be a misaligned 0x204 (on the host there was no need to add117 * padding between b and p, so p comes exactly 4 bytes after b). If the device118 * kernel tries to access s1.p, a misaligned address error occurs (as reported119 * by the CUDA plugin). By padding the begin address down to a multiple of 8 and120 * extending the size of the allocated chuck accordingly, the chuck on the121 * device will start at 0x200 with the padding (4 bytes), then &s1.b=0x204 and122 * &s1.p=0x208, as they should be to satisfy the alignment requirements.123 */124static const int64_t MaxAlignment = 16;125 126/// Return the alignment requirement of partially mapped structs, see127/// MaxAlignment above.128static uint64_t getPartialStructRequiredAlignment(void *HstPtrBase) {129 int LowestOneBit = __builtin_ffsl(reinterpret_cast<uintptr_t>(HstPtrBase));130 uint64_t BaseAlignment = 1 << (LowestOneBit - 1);131 return MaxAlignment < BaseAlignment ? MaxAlignment : BaseAlignment;132}133 134void handleTargetOutcome(bool Success, ident_t *Loc) {135 switch (OffloadPolicy::get(*PM).Kind) {136 case OffloadPolicy::DISABLED:137 if (Success) {138 FATAL_MESSAGE0(1, "expected no offloading while offloading is disabled");139 }140 break;141 case OffloadPolicy::MANDATORY:142 if (!Success) {143 if (getInfoLevel() & OMP_INFOTYPE_DUMP_TABLE) {144 auto ExclusiveDevicesAccessor = PM->getExclusiveDevicesAccessor();145 for (auto &Device : PM->devices(ExclusiveDevicesAccessor))146 dumpTargetPointerMappings(Loc, Device);147 } else148 FAILURE_MESSAGE("Consult https://openmp.llvm.org/design/Runtimes.html "149 "for debugging options.\n");150 151 if (!PM->getNumActivePlugins()) {152 FAILURE_MESSAGE(153 "No images found compatible with the installed hardware. ");154 155 llvm::SmallVector<llvm::StringRef> Archs;156 for (auto &Image : PM->deviceImages()) {157 const char *Start = reinterpret_cast<const char *>(158 Image.getExecutableImage().ImageStart);159 uint64_t Length =160 utils::getPtrDiff(Start, Image.getExecutableImage().ImageEnd);161 llvm::MemoryBufferRef Buffer(llvm::StringRef(Start, Length),162 /*Identifier=*/"");163 164 auto ObjectOrErr = llvm::object::ObjectFile::createObjectFile(Buffer);165 if (auto Err = ObjectOrErr.takeError()) {166 llvm::consumeError(std::move(Err));167 continue;168 }169 170 if (auto CPU = (*ObjectOrErr)->tryGetCPUName())171 Archs.push_back(*CPU);172 }173 fprintf(stderr, "Found %zu image(s): (%s)\n", Archs.size(),174 llvm::join(Archs, ",").c_str());175 }176 177 SourceInfo Info(Loc);178 if (Info.isAvailible())179 fprintf(stderr, "%s:%d:%d: ", Info.getFilename(), Info.getLine(),180 Info.getColumn());181 else182 FAILURE_MESSAGE("Source location information not present. Compile with "183 "-g or -gline-tables-only.\n");184 FATAL_MESSAGE0(185 1, "failure of target construct while offloading is mandatory");186 } else {187 if (getInfoLevel() & OMP_INFOTYPE_DUMP_TABLE) {188 auto ExclusiveDevicesAccessor = PM->getExclusiveDevicesAccessor();189 for (auto &Device : PM->devices(ExclusiveDevicesAccessor))190 dumpTargetPointerMappings(Loc, Device);191 }192 }193 break;194 }195}196 197static int32_t getParentIndex(int64_t Type) {198 return ((Type & OMP_TGT_MAPTYPE_MEMBER_OF) >> 48) - 1;199}200 201void *targetAllocExplicit(size_t Size, int DeviceNum, int Kind,202 const char *Name) {203 DP("Call to %s for device %d requesting %zu bytes\n", Name, DeviceNum, Size);204 205 if (Size <= 0) {206 DP("Call to %s with non-positive length\n", Name);207 return NULL;208 }209 210 void *Rc = NULL;211 212 if (DeviceNum == omp_get_initial_device()) {213 Rc = malloc(Size);214 DP("%s returns host ptr " DPxMOD "\n", Name, DPxPTR(Rc));215 return Rc;216 }217 218 auto DeviceOrErr = PM->getDevice(DeviceNum);219 if (!DeviceOrErr)220 FATAL_MESSAGE(DeviceNum, "%s", toString(DeviceOrErr.takeError()).c_str());221 222 Rc = DeviceOrErr->allocData(Size, nullptr, Kind);223 DP("%s returns device ptr " DPxMOD "\n", Name, DPxPTR(Rc));224 return Rc;225}226 227void targetFreeExplicit(void *DevicePtr, int DeviceNum, int Kind,228 const char *Name) {229 DP("Call to %s for device %d and address " DPxMOD "\n", Name, DeviceNum,230 DPxPTR(DevicePtr));231 232 if (!DevicePtr) {233 DP("Call to %s with NULL ptr\n", Name);234 return;235 }236 237 if (DeviceNum == omp_get_initial_device()) {238 free(DevicePtr);239 DP("%s deallocated host ptr\n", Name);240 return;241 }242 243 auto DeviceOrErr = PM->getDevice(DeviceNum);244 if (!DeviceOrErr)245 FATAL_MESSAGE(DeviceNum, "%s", toString(DeviceOrErr.takeError()).c_str());246 247 if (DeviceOrErr->deleteData(DevicePtr, Kind) == OFFLOAD_FAIL)248 FATAL_MESSAGE(DeviceNum, "%s",249 "Failed to deallocate device ptr. Set "250 "OFFLOAD_TRACK_ALLOCATION_TRACES=1 to track allocations.");251 252 DP("omp_target_free deallocated device ptr\n");253}254 255void *targetLockExplicit(void *HostPtr, size_t Size, int DeviceNum,256 const char *Name) {257 DP("Call to %s for device %d locking %zu bytes\n", Name, DeviceNum, Size);258 259 if (Size <= 0) {260 DP("Call to %s with non-positive length\n", Name);261 return NULL;262 }263 264 void *RC = NULL;265 266 auto DeviceOrErr = PM->getDevice(DeviceNum);267 if (!DeviceOrErr)268 FATAL_MESSAGE(DeviceNum, "%s", toString(DeviceOrErr.takeError()).c_str());269 270 int32_t Err = 0;271 Err = DeviceOrErr->RTL->data_lock(DeviceNum, HostPtr, Size, &RC);272 if (Err) {273 DP("Could not lock ptr %p\n", HostPtr);274 return nullptr;275 }276 DP("%s returns device ptr " DPxMOD "\n", Name, DPxPTR(RC));277 return RC;278}279 280void targetUnlockExplicit(void *HostPtr, int DeviceNum, const char *Name) {281 DP("Call to %s for device %d unlocking\n", Name, DeviceNum);282 283 auto DeviceOrErr = PM->getDevice(DeviceNum);284 if (!DeviceOrErr)285 FATAL_MESSAGE(DeviceNum, "%s", toString(DeviceOrErr.takeError()).c_str());286 287 DeviceOrErr->RTL->data_unlock(DeviceNum, HostPtr);288 DP("%s returns\n", Name);289}290 291/// Call the user-defined mapper function followed by the appropriate292// targetData* function (targetData{Begin,End,Update}).293int targetDataMapper(ident_t *Loc, DeviceTy &Device, void *ArgBase, void *Arg,294 int64_t ArgSize, int64_t ArgType, map_var_info_t ArgNames,295 void *ArgMapper, AsyncInfoTy &AsyncInfo,296 TargetDataFuncPtrTy TargetDataFunction,297 AttachInfoTy *AttachInfo = nullptr) {298 DP("Calling the mapper function " DPxMOD "\n", DPxPTR(ArgMapper));299 300 // The mapper function fills up Components.301 MapperComponentsTy MapperComponents;302 MapperFuncPtrTy MapperFuncPtr = (MapperFuncPtrTy)(ArgMapper);303 (*MapperFuncPtr)((void *)&MapperComponents, ArgBase, Arg, ArgSize, ArgType,304 ArgNames);305 306 // Construct new arrays for args_base, args, arg_sizes and arg_types307 // using the information in MapperComponents and call the corresponding308 // targetData* function using these new arrays.309 SmallVector<void *> MapperArgsBase(MapperComponents.Components.size());310 SmallVector<void *> MapperArgs(MapperComponents.Components.size());311 SmallVector<int64_t> MapperArgSizes(MapperComponents.Components.size());312 SmallVector<int64_t> MapperArgTypes(MapperComponents.Components.size());313 SmallVector<void *> MapperArgNames(MapperComponents.Components.size());314 315 for (unsigned I = 0, E = MapperComponents.Components.size(); I < E; ++I) {316 auto &C = MapperComponents.Components[I];317 MapperArgsBase[I] = C.Base;318 MapperArgs[I] = C.Begin;319 MapperArgSizes[I] = C.Size;320 MapperArgTypes[I] = C.Type;321 MapperArgNames[I] = C.Name;322 }323 324 int Rc = TargetDataFunction(Loc, Device, MapperComponents.Components.size(),325 MapperArgsBase.data(), MapperArgs.data(),326 MapperArgSizes.data(), MapperArgTypes.data(),327 MapperArgNames.data(), /*arg_mappers*/ nullptr,328 AsyncInfo, AttachInfo, /*FromMapper=*/true);329 330 return Rc;331}332 333/// Returns a buffer of the requested \p Size, to be used as the source for334/// `submitData`.335///336/// For small buffers (`Size <= sizeof(void*)`), uses \p AsyncInfo's337/// getVoidPtrLocation().338/// For larger buffers, creates a dynamic buffer which will be eventually339/// deleted by \p AsyncInfo's post-processing callback.340static char *getOrCreateSourceBufferForSubmitData(AsyncInfoTy &AsyncInfo,341 int64_t Size) {342 constexpr int64_t VoidPtrSize = sizeof(void *);343 344 if (Size <= VoidPtrSize) {345 void *&BufferElement = AsyncInfo.getVoidPtrLocation();346 return reinterpret_cast<char *>(&BufferElement);347 }348 349 // Create a dynamic buffer for larger data and schedule its deletion.350 char *DataBuffer = new char[Size];351 AsyncInfo.addPostProcessingFunction([DataBuffer]() {352 delete[] DataBuffer;353 return OFFLOAD_SUCCESS;354 });355 return DataBuffer;356}357 358/// Calculates the target pointee base by applying the host359/// pointee begin/base delta to the target pointee begin.360///361/// ```362/// TgtPteeBase = TgtPteeBegin - (HstPteeBegin - HstPteeBase)363/// ```364static void *calculateTargetPointeeBase(void *HstPteeBase, void *HstPteeBegin,365 void *TgtPteeBegin) {366 uint64_t Delta = reinterpret_cast<uint64_t>(HstPteeBegin) -367 reinterpret_cast<uint64_t>(HstPteeBase);368 void *TgtPteeBase = reinterpret_cast<void *>(369 reinterpret_cast<uint64_t>(TgtPteeBegin) - Delta);370 371 DP("HstPteeBase: " DPxMOD ", HstPteeBegin: " DPxMOD372 ", Delta (HstPteeBegin - HstPteeBase): %" PRIu64 ".\n",373 DPxPTR(HstPteeBase), DPxPTR(HstPteeBegin), Delta);374 DP("TgtPteeBase (TgtPteeBegin - Delta): " DPxMOD ", TgtPteeBegin : " DPxMOD375 "\n",376 DPxPTR(TgtPteeBase), DPxPTR(TgtPteeBegin));377 378 return TgtPteeBase;379}380 381/// Utility function to perform a pointer attachment operation.382///383/// For something like:384/// ```cpp385/// int *p;386/// ...387/// #pragma omp target enter data map(to:p[10:10])388/// ```389///390/// for which the attachment operation gets represented using:391/// ```392/// &p, &p[10], sizeof(p), ATTACH393/// ```394///395/// (Hst|Tgt)PtrAddr represents &p396/// (Hst|Tgt)PteeBase represents &p[0]397/// (Hst|Tgt)PteeBegin represents &p[10]398///399/// This function first computes the expected TgtPteeBase using:400/// `<Select>TgtPteeBase = TgtPteeBegin - (HstPteeBegin - HstPteeBase)`401///402/// and then attaches TgtPteeBase to TgtPtrAddr.403///404/// \p HstPtrSize represents the size of the pointer p. For C/C++, this405/// should be same as "sizeof(void*)" (say 8).406///407/// However, for Fortran, pointers/allocatables, which are also eligible for408/// "pointer-attachment", may be implemented using descriptors that contain the409/// address of the pointee in the first 8 bytes, but also contain other410/// information such as lower-bound/upper-bound etc in their subsequent fields.411///412/// For example, for the following:413/// ```fortran414/// integer, allocatable :: x(:)415/// integer, pointer :: p(:)416/// ...417/// p => x(10: 19)418/// ...419/// !$omp target enter data map(to:p(:))420/// ```421///422/// The map should trigger a pointer-attachment (assuming the pointer-attachment423/// conditions as noted on processAttachEntries are met) between the descriptor424/// for p, and its pointee data.425///426/// Since only the first 8 bytes of the descriptor contain the address of the427/// pointee, an attachment operation on device descriptors involves:428/// * Setting the first 8 bytes of the device descriptor to point the device429/// address of the pointee.430/// * Copying the remaining information about bounds/offset etc. from the host431/// descriptor to the device descriptor.432///433/// The function also handles pointer-attachment portion of PTR_AND_OBJ maps,434/// like:435/// ```436/// &p, &p[10], 10 * sizeof(p[10]), PTR_AND_OBJ437/// ```438/// by using `sizeof(void*)` as \p HstPtrSize.439static int performPointerAttachment(DeviceTy &Device, AsyncInfoTy &AsyncInfo,440 void **HstPtrAddr, void *HstPteeBase,441 void *HstPteeBegin, void **TgtPtrAddr,442 void *TgtPteeBegin, int64_t HstPtrSize,443 TargetPointerResultTy &PtrTPR) {444 assert(PtrTPR.getEntry() &&445 "Need a valid pointer entry to perform pointer-attachment");446 447 constexpr int64_t VoidPtrSize = sizeof(void *);448 assert(HstPtrSize >= VoidPtrSize && "PointerSize is too small");449 450 void *TgtPteeBase =451 calculateTargetPointeeBase(HstPteeBase, HstPteeBegin, TgtPteeBegin);452 453 // Add shadow pointer tracking454 if (!PtrTPR.getEntry()->addShadowPointer(455 ShadowPtrInfoTy{HstPtrAddr, TgtPtrAddr, TgtPteeBase, HstPtrSize})) {456 DP("Pointer " DPxMOD " is already attached to " DPxMOD "\n",457 DPxPTR(TgtPtrAddr), DPxPTR(TgtPteeBase));458 return OFFLOAD_SUCCESS;459 }460 461 DP("Update pointer (" DPxMOD ") -> [" DPxMOD "]\n", DPxPTR(TgtPtrAddr),462 DPxPTR(TgtPteeBase));463 464 // Lambda to handle submitData result and perform final steps.465 auto HandleSubmitResult = [&](int SubmitResult) -> int {466 if (SubmitResult != OFFLOAD_SUCCESS) {467 REPORT("Failed to update pointer on device.\n");468 return OFFLOAD_FAIL;469 }470 471 if (PtrTPR.getEntry()->addEventIfNecessary(Device, AsyncInfo) !=472 OFFLOAD_SUCCESS)473 return OFFLOAD_FAIL;474 475 return OFFLOAD_SUCCESS;476 };477 478 // Get a buffer to be used as the source for data submission.479 char *SrcBuffer = getOrCreateSourceBufferForSubmitData(AsyncInfo, HstPtrSize);480 481 // The pointee's address should occupy the first VoidPtrSize bytes482 // irrespective of HstPtrSize.483 std::memcpy(SrcBuffer, &TgtPteeBase, VoidPtrSize);484 485 // For larger "pointers" (e.g., Fortran descriptors), copy remaining486 // descriptor fields from the host descriptor into the buffer.487 if (HstPtrSize > VoidPtrSize) {488 uint64_t HstDescriptorFieldsSize = HstPtrSize - VoidPtrSize;489 void *HstDescriptorFieldsAddr =490 reinterpret_cast<char *>(HstPtrAddr) + VoidPtrSize;491 std::memcpy(SrcBuffer + VoidPtrSize, HstDescriptorFieldsAddr,492 HstDescriptorFieldsSize);493 494 DP("Updating %" PRId64 " bytes of descriptor (" DPxMOD495 ") (pointer + %" PRId64 " additional bytes from host descriptor " DPxMOD496 ")\n",497 HstPtrSize, DPxPTR(TgtPtrAddr), HstDescriptorFieldsSize,498 DPxPTR(HstDescriptorFieldsAddr));499 }500 501 // Submit the populated source buffer to device.502 int SubmitResult = Device.submitData(TgtPtrAddr, SrcBuffer, HstPtrSize,503 AsyncInfo, PtrTPR.getEntry());504 return HandleSubmitResult(SubmitResult);505}506 507/// Internal function to do the mapping and transfer the data to the device508int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,509 void **ArgsBase, void **Args, int64_t *ArgSizes,510 int64_t *ArgTypes, map_var_info_t *ArgNames,511 void **ArgMappers, AsyncInfoTy &AsyncInfo,512 AttachInfoTy *AttachInfo, bool FromMapper) {513 assert(AttachInfo && "AttachInfo must be available for targetDataBegin for "514 "handling ATTACH map-types.");515 // process each input.516 for (int32_t I = 0; I < ArgNum; ++I) {517 // Ignore private variables and arrays - there is no mapping for them.518 if ((ArgTypes[I] & OMP_TGT_MAPTYPE_LITERAL) ||519 (ArgTypes[I] & OMP_TGT_MAPTYPE_PRIVATE))520 continue;521 TIMESCOPE_WITH_DETAILS_AND_IDENT(522 "HostToDev", "Size=" + std::to_string(ArgSizes[I]) + "B", Loc);523 if (ArgMappers && ArgMappers[I]) {524 // Instead of executing the regular path of targetDataBegin, call the525 // targetDataMapper variant which will call targetDataBegin again526 // with new arguments.527 DP("Calling targetDataMapper for the %dth argument\n", I);528 529 map_var_info_t ArgName = (!ArgNames) ? nullptr : ArgNames[I];530 int Rc = targetDataMapper(Loc, Device, ArgsBase[I], Args[I], ArgSizes[I],531 ArgTypes[I], ArgName, ArgMappers[I], AsyncInfo,532 targetDataBegin, AttachInfo);533 534 if (Rc != OFFLOAD_SUCCESS) {535 REPORT("Call to targetDataBegin via targetDataMapper for custom mapper"536 " failed.\n");537 return OFFLOAD_FAIL;538 }539 540 // Skip the rest of this function, continue to the next argument.541 continue;542 }543 544 void *HstPtrBegin = Args[I];545 void *HstPtrBase = ArgsBase[I];546 int64_t DataSize = ArgSizes[I];547 map_var_info_t HstPtrName = (!ArgNames) ? nullptr : ArgNames[I];548 549 // ATTACH map-types are supposed to be handled after all mapping for the550 // construct is done. Defer their processing.551 if (ArgTypes[I] & OMP_TGT_MAPTYPE_ATTACH) {552 const bool IsCorrespondingPointerInit =553 (ArgTypes[I] & OMP_TGT_MAPTYPE_PRIVATE);554 // We don't need to keep track of PRIVATE | ATTACH entries. They555 // represent corresponding-pointer-initialization, and are handled556 // similar to firstprivate (PRIVATE | TO) entries by557 // PrivateArgumentManager.558 if (!IsCorrespondingPointerInit)559 AttachInfo->AttachEntries.emplace_back(560 /*PointerBase=*/HstPtrBase, /*PointeeBegin=*/HstPtrBegin,561 /*PointerSize=*/DataSize, /*MapType=*/ArgTypes[I],562 /*PointeeName=*/HstPtrName);563 564 DP("Deferring ATTACH map-type processing for argument %d\n", I);565 continue;566 }567 568 // Adjust for proper alignment if this is a combined entry (for structs).569 // Look at the next argument - if that is MEMBER_OF this one, then this one570 // is a combined entry.571 int64_t TgtPadding = 0;572 const int NextI = I + 1;573 if (getParentIndex(ArgTypes[I]) < 0 && NextI < ArgNum &&574 getParentIndex(ArgTypes[NextI]) == I) {575 int64_t Alignment = getPartialStructRequiredAlignment(HstPtrBase);576 TgtPadding = (int64_t)HstPtrBegin % Alignment;577 if (TgtPadding) {578 DP("Using a padding of %" PRId64 " bytes for begin address " DPxMOD579 "\n",580 TgtPadding, DPxPTR(HstPtrBegin));581 }582 }583 584 // Address of pointer on the host and device, respectively.585 void *PointerHstPtrBegin, *PointerTgtPtrBegin;586 TargetPointerResultTy PointerTpr;587 bool IsHostPtr = false;588 bool IsImplicit = ArgTypes[I] & OMP_TGT_MAPTYPE_IMPLICIT;589 // Force the creation of a device side copy of the data when:590 // a close map modifier was associated with a map that contained a to.591 bool HasCloseModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_CLOSE;592 bool HasPresentModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_PRESENT;593 bool HasHoldModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_OMPX_HOLD;594 // UpdateRef is based on MEMBER_OF instead of TARGET_PARAM because if we595 // have reached this point via __tgt_target_data_begin and not __tgt_target596 // then no argument is marked as TARGET_PARAM ("omp target data map" is not597 // associated with a target region, so there are no target parameters). This598 // may be considered a hack, we could revise the scheme in the future.599 bool UpdateRef =600 !(ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF) && !(FromMapper && I == 0);601 602 MappingInfoTy::HDTTMapAccessorTy HDTTMap =603 Device.getMappingInfo().HostDataToTargetMap.getExclusiveAccessor();604 if (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ) {605 DP("Has a pointer entry: \n");606 // Base is address of pointer.607 //608 // Usually, the pointer is already allocated by this time. For example:609 //610 // #pragma omp target map(s.p[0:N])611 //612 // The map entry for s comes first, and the PTR_AND_OBJ entry comes613 // afterward, so the pointer is already allocated by the time the614 // PTR_AND_OBJ entry is handled below, and PointerTgtPtrBegin is thus615 // non-null. However, "declare target link" can produce a PTR_AND_OBJ616 // entry for a global that might not already be allocated by the time the617 // PTR_AND_OBJ entry is handled below, and so the allocation might fail618 // when HasPresentModifier.619 PointerTpr = Device.getMappingInfo().getTargetPointer(620 HDTTMap, HstPtrBase, HstPtrBase, /*TgtPadding=*/0, sizeof(void *),621 /*HstPtrName=*/nullptr,622 /*HasFlagTo=*/false, /*HasFlagAlways=*/false, IsImplicit, UpdateRef,623 HasCloseModifier, HasPresentModifier, HasHoldModifier, AsyncInfo,624 /*OwnedTPR=*/nullptr, /*ReleaseHDTTMap=*/false);625 PointerTgtPtrBegin = PointerTpr.TargetPointer;626 IsHostPtr = PointerTpr.Flags.IsHostPointer;627 if (!PointerTgtPtrBegin) {628 REPORT("Call to getTargetPointer returned null pointer (%s).\n",629 HasPresentModifier ? "'present' map type modifier"630 : "device failure or illegal mapping");631 return OFFLOAD_FAIL;632 }633 634 // Track new allocation, for eventual use in attachment decision-making.635 if (PointerTpr.Flags.IsNewEntry && !IsHostPtr)636 AttachInfo->NewAllocations[HstPtrBase] = sizeof(void *);637 638 DP("There are %zu bytes allocated at target address " DPxMOD " - is%s new"639 "\n",640 sizeof(void *), DPxPTR(PointerTgtPtrBegin),641 (PointerTpr.Flags.IsNewEntry ? "" : " not"));642 PointerHstPtrBegin = HstPtrBase;643 // modify current entry.644 HstPtrBase = *reinterpret_cast<void **>(HstPtrBase);645 // No need to update pointee ref count for the first element of the646 // subelement that comes from mapper.647 UpdateRef =648 (!FromMapper || I != 0); // subsequently update ref count of pointee649 }650 651 const bool HasFlagTo = ArgTypes[I] & OMP_TGT_MAPTYPE_TO;652 const bool HasFlagAlways = ArgTypes[I] & OMP_TGT_MAPTYPE_ALWAYS;653 // Note that HDTTMap will be released in getTargetPointer.654 auto TPR = Device.getMappingInfo().getTargetPointer(655 HDTTMap, HstPtrBegin, HstPtrBase, TgtPadding, DataSize, HstPtrName,656 HasFlagTo, HasFlagAlways, IsImplicit, UpdateRef, HasCloseModifier,657 HasPresentModifier, HasHoldModifier, AsyncInfo, PointerTpr.getEntry());658 void *TgtPtrBegin = TPR.TargetPointer;659 IsHostPtr = TPR.Flags.IsHostPointer;660 // If data_size==0, then the argument could be a zero-length pointer to661 // NULL, so getOrAlloc() returning NULL is not an error.662 if (!TgtPtrBegin && (DataSize || HasPresentModifier)) {663 REPORT("Call to getTargetPointer returned null pointer (%s).\n",664 HasPresentModifier ? "'present' map type modifier"665 : "device failure or illegal mapping");666 return OFFLOAD_FAIL;667 }668 669 // Track new allocation, for eventual use in attachment decision-making.670 if (TPR.Flags.IsNewEntry && !IsHostPtr && TgtPtrBegin)671 AttachInfo->NewAllocations[HstPtrBegin] = DataSize;672 673 DP("There are %" PRId64 " bytes allocated at target address " DPxMOD674 " - is%s new\n",675 DataSize, DPxPTR(TgtPtrBegin), (TPR.Flags.IsNewEntry ? "" : " not"));676 677 if (ArgTypes[I] & OMP_TGT_MAPTYPE_RETURN_PARAM) {678 uintptr_t Delta = (uintptr_t)HstPtrBegin - (uintptr_t)HstPtrBase;679 void *TgtPtrBase = (void *)((uintptr_t)TgtPtrBegin - Delta);680 DP("Returning device pointer " DPxMOD "\n", DPxPTR(TgtPtrBase));681 ArgsBase[I] = TgtPtrBase;682 }683 684 if (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ && !IsHostPtr) {685 int Ret = performPointerAttachment(686 Device, AsyncInfo, reinterpret_cast<void **>(PointerHstPtrBegin),687 HstPtrBase, HstPtrBegin,688 reinterpret_cast<void **>(PointerTgtPtrBegin), TgtPtrBegin,689 sizeof(void *), PointerTpr);690 if (Ret != OFFLOAD_SUCCESS)691 return OFFLOAD_FAIL;692 }693 694 // Check if variable can be used on the device:695 bool IsStructMember = ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF;696 if (getInfoLevel() & OMP_INFOTYPE_EMPTY_MAPPING && ArgTypes[I] != 0 &&697 !IsStructMember && !IsImplicit && !TPR.isPresent() &&698 !TPR.isContained() && !TPR.isHostPointer())699 INFO(OMP_INFOTYPE_EMPTY_MAPPING, Device.DeviceID,700 "variable %s does not have a valid device counterpart\n",701 (HstPtrName) ? getNameFromMapping(HstPtrName).c_str() : "unknown");702 }703 704 return OFFLOAD_SUCCESS;705}706 707/// Process deferred ATTACH map entries collected during targetDataBegin.708///709/// From OpenMP's perspective, when mapping something that has a base pointer,710/// such as:711/// ```cpp712/// int *p;713/// #pragma omp enter target data map(to: p[10:20])714/// ```715///716/// a pointer-attachment between p and &p[10] should occur if both p and717/// p[10] are present on the device after doing all allocations for all maps718/// on the construct, and one of the following is true:719///720/// * The pointer p was newly allocated while handling the construct721/// * The pointee p[10:20] was newly allocated while handling the construct722/// * attach(always) map-type modifier was specified (OpenMP 6.1)723///724/// That's why we collect all attach entries and new memory allocations during725/// targetDataBegin, and use that information to make the decision of whether726/// to perform a pointer-attachment or not here, after maps have been handled.727///728/// Additionally, once we decide that a pointer-attachment should be performed,729/// we need to make sure that it happens after any previously submitted data730/// transfers have completed, to avoid the possibility of the pending transfers731/// clobbering the attachment. For example:732///733/// ```cpp734/// int *p = ...;735/// int **pp = &p;736/// map(to: pp[0], p[0])737/// ```738///739/// Which would be represented by:740/// ```741/// &pp[0], &pp[0], sizeof(pp[0]), TO (1)742/// &p[0], &p[0], sizeof(p[0]), TO (2)743///744/// &pp, &pp[0], sizeof(pp), ATTACH (3)745/// &p, &p[0], sizeof(p), ATTACH (4)746/// ```747///748/// (4) and (1) are both trying to modify the device memory corresponding to749/// `&p`. So, if we decide that (4) should do an attachment, we also need to750/// ensure that (4) happens after (1) is complete.751///752/// For this purpose, we insert a data_fence before the first753/// pointer-attachment, (3), to ensure that all pending transfers finish first.754int processAttachEntries(DeviceTy &Device, AttachInfoTy &AttachInfo,755 AsyncInfoTy &AsyncInfo) {756 // Report all tracked allocations from both main loop and ATTACH processing757 if (!AttachInfo.NewAllocations.empty()) {758 DP("Tracked %u total new allocations:\n",759 (unsigned)AttachInfo.NewAllocations.size());760 for ([[maybe_unused]] const auto &Alloc : AttachInfo.NewAllocations) {761 DP(" Host ptr: " DPxMOD ", Size: %" PRId64 " bytes\n",762 DPxPTR(Alloc.first), Alloc.second);763 }764 }765 766 if (AttachInfo.AttachEntries.empty())767 return OFFLOAD_SUCCESS;768 769 DP("Processing %zu deferred ATTACH map entries\n",770 AttachInfo.AttachEntries.size());771 772 int Ret = OFFLOAD_SUCCESS;773 bool IsFirstPointerAttachment = true;774 for (size_t EntryIdx = 0; EntryIdx < AttachInfo.AttachEntries.size();775 ++EntryIdx) {776 const auto &AttachEntry = AttachInfo.AttachEntries[EntryIdx];777 778 void **HstPtr = reinterpret_cast<void **>(AttachEntry.PointerBase);779 780 void *HstPteeBase = *HstPtr;781 void *HstPteeBegin = AttachEntry.PointeeBegin;782 783 int64_t PtrSize = AttachEntry.PointerSize;784 int64_t MapType = AttachEntry.MapType;785 786 DP("Processing ATTACH entry %zu: HstPtr=" DPxMOD ", HstPteeBegin=" DPxMOD787 ", Size=%" PRId64 ", Type=0x%" PRIx64 "\n",788 EntryIdx, DPxPTR(HstPtr), DPxPTR(HstPteeBegin), PtrSize, MapType);789 790 const bool IsAttachAlways = MapType & OMP_TGT_MAPTYPE_ALWAYS;791 792 // Lambda to check if a pointer was newly allocated793 auto WasNewlyAllocated = [&](void *Ptr, const char *PtrName) {794 bool IsNewlyAllocated =795 llvm::any_of(AttachInfo.NewAllocations, [&](const auto &Alloc) {796 void *AllocPtr = Alloc.first;797 int64_t AllocSize = Alloc.second;798 return Ptr >= AllocPtr &&799 Ptr < reinterpret_cast<void *>(800 reinterpret_cast<char *>(AllocPtr) + AllocSize);801 });802 DP("Attach %s " DPxMOD " was newly allocated: %s\n", PtrName, DPxPTR(Ptr),803 IsNewlyAllocated ? "yes" : "no");804 return IsNewlyAllocated;805 };806 807 // Only process ATTACH if either the pointee or the pointer was newly808 // allocated, or the ALWAYS flag is set.809 if (!IsAttachAlways && !WasNewlyAllocated(HstPteeBegin, "pointee") &&810 !WasNewlyAllocated(HstPtr, "pointer")) {811 DP("Skipping ATTACH entry %zu: neither pointer nor pointee was newly "812 "allocated and no ALWAYS flag\n",813 EntryIdx);814 continue;815 }816 817 // Lambda to perform target pointer lookup and validation818 auto LookupTargetPointer =819 [&](void *Ptr, int64_t Size,820 const char *PtrType) -> std::optional<TargetPointerResultTy> {821 // ATTACH map-type does not change ref-count, or do any allocation822 // We just need to do a lookup for the pointer/pointee.823 TargetPointerResultTy TPR = Device.getMappingInfo().getTgtPtrBegin(824 Ptr, Size, /*UpdateRefCount=*/false,825 /*UseHoldRefCount=*/false, /*MustContain=*/true);826 827 DP("Attach %s lookup - IsPresent=%s, IsHostPtr=%s\n", PtrType,828 TPR.isPresent() ? "yes" : "no",829 TPR.Flags.IsHostPointer ? "yes" : "no");830 831 if (!TPR.isPresent()) {832 DP("Skipping ATTACH entry %zu: %s not present on device\n", EntryIdx,833 PtrType);834 return std::nullopt;835 }836 if (TPR.Flags.IsHostPointer) {837 DP("Skipping ATTACH entry %zu: device version of the %s is a host "838 "pointer.\n",839 EntryIdx, PtrType);840 return std::nullopt;841 }842 843 return TPR;844 };845 846 // Get device version of the pointee (e.g., &p[10]) first, as we can847 // release its TPR after extracting the pointer value.848 void *TgtPteeBegin = [&]() -> void * {849 if (auto PteeTPROpt = LookupTargetPointer(HstPteeBegin, 0, "pointee"))850 return PteeTPROpt->TargetPointer;851 return nullptr;852 }();853 854 if (!TgtPteeBegin)855 continue;856 857 // Get device version of the pointer (e.g., &p) next. We need to keep its858 // TPR for use in shadow-pointer handling during pointer-attachment.859 auto PtrTPROpt = LookupTargetPointer(HstPtr, PtrSize, "pointer");860 if (!PtrTPROpt)861 continue;862 TargetPointerResultTy &PtrTPR = *PtrTPROpt;863 void **TgtPtrBase = reinterpret_cast<void **>(PtrTPR.TargetPointer);864 865 // Insert a data-fence before the first pointer-attachment.866 if (IsFirstPointerAttachment) {867 IsFirstPointerAttachment = false;868 DP("Inserting a data fence before the first pointer attachment.\n");869 Ret = Device.dataFence(AsyncInfo);870 if (Ret != OFFLOAD_SUCCESS) {871 REPORT("Failed to insert data fence.\n");872 return OFFLOAD_FAIL;873 }874 }875 876 // Do the pointer-attachment, i.e. update the device pointer to point to877 // device pointee.878 Ret = performPointerAttachment(Device, AsyncInfo, HstPtr, HstPteeBase,879 HstPteeBegin, TgtPtrBase, TgtPteeBegin,880 PtrSize, PtrTPR);881 if (Ret != OFFLOAD_SUCCESS)882 return OFFLOAD_FAIL;883 884 DP("ATTACH entry %zu processed successfully\n", EntryIdx);885 }886 887 return OFFLOAD_SUCCESS;888}889 890namespace {891/// This structure contains information to deallocate a target pointer, aka.892/// used to fix up the shadow map and potentially delete the entry from the893/// mapping table via \p DeviceTy::deallocTgtPtr.894struct PostProcessingInfo {895 /// Host pointer used to look up into the map table896 void *HstPtrBegin;897 898 /// Size of the data899 int64_t DataSize;900 901 /// The mapping type (bitfield).902 int64_t ArgType;903 904 /// The target pointer information.905 TargetPointerResultTy TPR;906 907 PostProcessingInfo(void *HstPtr, int64_t Size, int64_t ArgType,908 TargetPointerResultTy &&TPR)909 : HstPtrBegin(HstPtr), DataSize(Size), ArgType(ArgType),910 TPR(std::move(TPR)) {}911};912 913} // namespace914 915/// Applies the necessary post-processing procedures to entries listed in \p916/// EntriesInfo after the execution of all device side operations from a target917/// data end. This includes the update of pointers at the host and removal of918/// device buffer when needed. It returns OFFLOAD_FAIL or OFFLOAD_SUCCESS919/// according to the successfulness of the operations.920[[nodiscard]] static int921postProcessingTargetDataEnd(DeviceTy *Device,922 SmallVector<PostProcessingInfo> &EntriesInfo) {923 int Ret = OFFLOAD_SUCCESS;924 925 for (auto &[HstPtrBegin, DataSize, ArgType, TPR] : EntriesInfo) {926 bool DelEntry = !TPR.isHostPointer();927 928 // If the last element from the mapper (for end transfer args comes in929 // reverse order), do not remove the partial entry, the parent struct still930 // exists.931 if ((ArgType & OMP_TGT_MAPTYPE_MEMBER_OF) &&932 !(ArgType & OMP_TGT_MAPTYPE_PTR_AND_OBJ)) {933 DelEntry = false; // protect parent struct from being deallocated934 }935 936 // If we marked the entry to be deleted we need to verify no other937 // thread reused it by now. If deletion is still supposed to happen by938 // this thread LR will be set and exclusive access to the HDTT map939 // will avoid another thread reusing the entry now. Note that we do940 // not request (exclusive) access to the HDTT map if DelEntry is941 // not set.942 MappingInfoTy::HDTTMapAccessorTy HDTTMap =943 Device->getMappingInfo().HostDataToTargetMap.getExclusiveAccessor();944 945 // We cannot use a lock guard because we may end up delete the mutex.946 // We also explicitly unlocked the entry after it was put in the EntriesInfo947 // so it can be reused.948 TPR.getEntry()->lock();949 auto *Entry = TPR.getEntry();950 951 const bool IsNotLastUser = Entry->decDataEndThreadCount() != 0;952 if (DelEntry && (Entry->getTotalRefCount() != 0 || IsNotLastUser)) {953 // The thread is not in charge of deletion anymore. Give up access954 // to the HDTT map and unset the deletion flag.955 HDTTMap.destroy();956 DelEntry = false;957 }958 959 // If we copied back to the host a struct/array containing pointers, or960 // Fortran descriptors (which are larger than a "void *"), we need to961 // restore the original host pointer/descriptor values from their shadow962 // copies. If the struct is going to be deallocated, remove any remaining963 // shadow pointer entries for this struct.964 const bool HasFrom = ArgType & OMP_TGT_MAPTYPE_FROM;965 if (HasFrom) {966 Entry->foreachShadowPointerInfo([&](const ShadowPtrInfoTy &ShadowPtr) {967 constexpr int64_t VoidPtrSize = sizeof(void *);968 if (ShadowPtr.PtrSize > VoidPtrSize) {969 DP("Restoring host descriptor " DPxMOD970 " to its original content (%" PRId64971 " bytes), containing pointee address " DPxMOD "\n",972 DPxPTR(ShadowPtr.HstPtrAddr), ShadowPtr.PtrSize,973 DPxPTR(ShadowPtr.HstPtrContent.data()));974 } else {975 DP("Restoring host pointer " DPxMOD " to its original value " DPxMOD976 "\n",977 DPxPTR(ShadowPtr.HstPtrAddr),978 DPxPTR(ShadowPtr.HstPtrContent.data()));979 }980 std::memcpy(ShadowPtr.HstPtrAddr, ShadowPtr.HstPtrContent.data(),981 ShadowPtr.PtrSize);982 return OFFLOAD_SUCCESS;983 });984 }985 986 // Give up the lock as we either don't need it anymore (e.g., done with987 // TPR), or erase TPR.988 TPR.setEntry(nullptr);989 990 if (!DelEntry)991 continue;992 993 Ret = Device->getMappingInfo().eraseMapEntry(HDTTMap, Entry, DataSize);994 // Entry is already remove from the map, we can unlock it now.995 HDTTMap.destroy();996 Ret |= Device->getMappingInfo().deallocTgtPtrAndEntry(Entry, DataSize);997 if (Ret != OFFLOAD_SUCCESS) {998 REPORT("Deallocating data from device failed.\n");999 break;1000 }1001 }1002 1003 delete &EntriesInfo;1004 return Ret;1005}1006 1007/// Internal function to undo the mapping and retrieve the data from the device.1008int targetDataEnd(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,1009 void **ArgBases, void **Args, int64_t *ArgSizes,1010 int64_t *ArgTypes, map_var_info_t *ArgNames,1011 void **ArgMappers, AsyncInfoTy &AsyncInfo,1012 AttachInfoTy *AttachInfo, bool FromMapper) {1013 int Ret = OFFLOAD_SUCCESS;1014 auto *PostProcessingPtrs = new SmallVector<PostProcessingInfo>();1015 // process each input.1016 for (int32_t I = ArgNum - 1; I >= 0; --I) {1017 // Ignore private variables and arrays - there is no mapping for them.1018 // Also, ignore the use_device_ptr directive, it has no effect here.1019 if ((ArgTypes[I] & OMP_TGT_MAPTYPE_LITERAL) ||1020 (ArgTypes[I] & OMP_TGT_MAPTYPE_PRIVATE))1021 continue;1022 1023 // Ignore ATTACH entries - they should only be honored on map-entering1024 // directives. They may be encountered here while handling the "end" part of1025 // "#pragma omp target".1026 if (ArgTypes[I] & OMP_TGT_MAPTYPE_ATTACH) {1027 DP("Ignoring ATTACH entry %d in targetDataEnd\n", I);1028 continue;1029 }1030 1031 if (ArgMappers && ArgMappers[I]) {1032 // Instead of executing the regular path of targetDataEnd, call the1033 // targetDataMapper variant which will call targetDataEnd again1034 // with new arguments.1035 DP("Calling targetDataMapper for the %dth argument\n", I);1036 1037 map_var_info_t ArgName = (!ArgNames) ? nullptr : ArgNames[I];1038 Ret = targetDataMapper(Loc, Device, ArgBases[I], Args[I], ArgSizes[I],1039 ArgTypes[I], ArgName, ArgMappers[I], AsyncInfo,1040 targetDataEnd);1041 1042 if (Ret != OFFLOAD_SUCCESS) {1043 REPORT("Call to targetDataEnd via targetDataMapper for custom mapper"1044 " failed.\n");1045 return OFFLOAD_FAIL;1046 }1047 1048 // Skip the rest of this function, continue to the next argument.1049 continue;1050 }1051 1052 void *HstPtrBegin = Args[I];1053 int64_t DataSize = ArgSizes[I];1054 bool IsImplicit = ArgTypes[I] & OMP_TGT_MAPTYPE_IMPLICIT;1055 bool UpdateRef = (!(ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF) ||1056 (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ)) &&1057 !(FromMapper && I == 0);1058 bool ForceDelete = ArgTypes[I] & OMP_TGT_MAPTYPE_DELETE;1059 bool HasPresentModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_PRESENT;1060 bool HasHoldModifier = ArgTypes[I] & OMP_TGT_MAPTYPE_OMPX_HOLD;1061 1062 // If PTR_AND_OBJ, HstPtrBegin is address of pointee1063 TargetPointerResultTy TPR = Device.getMappingInfo().getTgtPtrBegin(1064 HstPtrBegin, DataSize, UpdateRef, HasHoldModifier, !IsImplicit,1065 ForceDelete, /*FromDataEnd=*/true);1066 void *TgtPtrBegin = TPR.TargetPointer;1067 if (!TPR.isPresent() && !TPR.isHostPointer() &&1068 (DataSize || HasPresentModifier)) {1069 DP("Mapping does not exist (%s)\n",1070 (HasPresentModifier ? "'present' map type modifier" : "ignored"));1071 if (HasPresentModifier) {1072 // OpenMP 5.1, sec. 2.21.7.1 "map Clause", p. 350 L10-13:1073 // "If a map clause appears on a target, target data, target enter data1074 // or target exit data construct with a present map-type-modifier then1075 // on entry to the region if the corresponding list item does not appear1076 // in the device data environment then an error occurs and the program1077 // terminates."1078 //1079 // This should be an error upon entering an "omp target exit data". It1080 // should not be an error upon exiting an "omp target data" or "omp1081 // target". For "omp target data", Clang thus doesn't include present1082 // modifiers for end calls. For "omp target", we have not found a valid1083 // OpenMP program for which the error matters: it appears that, if a1084 // program can guarantee that data is present at the beginning of an1085 // "omp target" region so that there's no error there, that data is also1086 // guaranteed to be present at the end.1087 MESSAGE("device mapping required by 'present' map type modifier does "1088 "not exist for host address " DPxMOD " (%" PRId64 " bytes)",1089 DPxPTR(HstPtrBegin), DataSize);1090 return OFFLOAD_FAIL;1091 }1092 } else {1093 DP("There are %" PRId64 " bytes allocated at target address " DPxMOD1094 " - is%s last\n",1095 DataSize, DPxPTR(TgtPtrBegin), (TPR.Flags.IsLast ? "" : " not"));1096 }1097 1098 // OpenMP 5.1, sec. 2.21.7.1 "map Clause", p. 351 L14-16:1099 // "If the map clause appears on a target, target data, or target exit data1100 // construct and a corresponding list item of the original list item is not1101 // present in the device data environment on exit from the region then the1102 // list item is ignored."1103 if (!TPR.isPresent())1104 continue;1105 1106 // Move data back to the host1107 const bool HasAlways = ArgTypes[I] & OMP_TGT_MAPTYPE_ALWAYS;1108 const bool HasFrom = ArgTypes[I] & OMP_TGT_MAPTYPE_FROM;1109 if (HasFrom && (HasAlways || TPR.Flags.IsLast) &&1110 !TPR.Flags.IsHostPointer && DataSize != 0) {1111 DP("Moving %" PRId64 " bytes (tgt:" DPxMOD ") -> (hst:" DPxMOD ")\n",1112 DataSize, DPxPTR(TgtPtrBegin), DPxPTR(HstPtrBegin));1113 TIMESCOPE_WITH_DETAILS_AND_IDENT(1114 "DevToHost", "Size=" + std::to_string(DataSize) + "B", Loc);1115 // Wait for any previous transfer if an event is present.1116 if (void *Event = TPR.getEntry()->getEvent()) {1117 if (Device.waitEvent(Event, AsyncInfo) != OFFLOAD_SUCCESS) {1118 REPORT("Failed to wait for event " DPxMOD ".\n", DPxPTR(Event));1119 return OFFLOAD_FAIL;1120 }1121 }1122 1123 Ret = Device.retrieveData(HstPtrBegin, TgtPtrBegin, DataSize, AsyncInfo,1124 TPR.getEntry());1125 if (Ret != OFFLOAD_SUCCESS) {1126 REPORT("Copying data from device failed.\n");1127 return OFFLOAD_FAIL;1128 }1129 1130 // As we are expecting to delete the entry the d2h copy might race1131 // with another one that also tries to delete the entry. This happens1132 // as the entry can be reused and the reuse might happen after the1133 // copy-back was issued but before it completed. Since the reuse might1134 // also copy-back a value we would race.1135 if (TPR.Flags.IsLast) {1136 if (TPR.getEntry()->addEventIfNecessary(Device, AsyncInfo) !=1137 OFFLOAD_SUCCESS)1138 return OFFLOAD_FAIL;1139 }1140 }1141 1142 // Add pointer to the buffer for post-synchronize processing.1143 PostProcessingPtrs->emplace_back(HstPtrBegin, DataSize, ArgTypes[I],1144 std::move(TPR));1145 PostProcessingPtrs->back().TPR.getEntry()->unlock();1146 }1147 1148 // Add post-processing functions1149 // TODO: We might want to remove `mutable` in the future by not changing the1150 // captured variables somehow.1151 AsyncInfo.addPostProcessingFunction([=, Device = &Device]() mutable -> int {1152 return postProcessingTargetDataEnd(Device, *PostProcessingPtrs);1153 });1154 1155 return Ret;1156}1157 1158static int targetDataContiguous(ident_t *Loc, DeviceTy &Device, void *ArgsBase,1159 void *HstPtrBegin, int64_t ArgSize,1160 int64_t ArgType, AsyncInfoTy &AsyncInfo) {1161 TargetPointerResultTy TPR = Device.getMappingInfo().getTgtPtrBegin(1162 HstPtrBegin, ArgSize, /*UpdateRefCount=*/false,1163 /*UseHoldRefCount=*/false, /*MustContain=*/true);1164 void *TgtPtrBegin = TPR.TargetPointer;1165 if (!TPR.isPresent()) {1166 DP("hst data:" DPxMOD " not found, becomes a noop\n", DPxPTR(HstPtrBegin));1167 if (ArgType & OMP_TGT_MAPTYPE_PRESENT) {1168 MESSAGE("device mapping required by 'present' motion modifier does not "1169 "exist for host address " DPxMOD " (%" PRId64 " bytes)",1170 DPxPTR(HstPtrBegin), ArgSize);1171 return OFFLOAD_FAIL;1172 }1173 return OFFLOAD_SUCCESS;1174 }1175 1176 if (TPR.Flags.IsHostPointer) {1177 DP("hst data:" DPxMOD " unified and shared, becomes a noop\n",1178 DPxPTR(HstPtrBegin));1179 return OFFLOAD_SUCCESS;1180 }1181 1182 if (ArgType & OMP_TGT_MAPTYPE_TO) {1183 DP("Moving %" PRId64 " bytes (hst:" DPxMOD ") -> (tgt:" DPxMOD ")\n",1184 ArgSize, DPxPTR(HstPtrBegin), DPxPTR(TgtPtrBegin));1185 int Ret = Device.submitData(TgtPtrBegin, HstPtrBegin, ArgSize, AsyncInfo,1186 TPR.getEntry());1187 if (Ret != OFFLOAD_SUCCESS) {1188 REPORT("Copying data to device failed.\n");1189 return OFFLOAD_FAIL;1190 }1191 if (TPR.getEntry()) {1192 int Ret = TPR.getEntry()->foreachShadowPointerInfo(1193 [&](ShadowPtrInfoTy &ShadowPtr) {1194 constexpr int64_t VoidPtrSize = sizeof(void *);1195 if (ShadowPtr.PtrSize > VoidPtrSize) {1196 DP("Restoring target descriptor " DPxMOD1197 " to its original content (%" PRId641198 " bytes), containing pointee address " DPxMOD "\n",1199 DPxPTR(ShadowPtr.TgtPtrAddr), ShadowPtr.PtrSize,1200 DPxPTR(ShadowPtr.TgtPtrContent.data()));1201 } else {1202 DP("Restoring target pointer " DPxMOD1203 " to its original value " DPxMOD "\n",1204 DPxPTR(ShadowPtr.TgtPtrAddr),1205 DPxPTR(ShadowPtr.TgtPtrContent.data()));1206 }1207 Ret = Device.submitData(ShadowPtr.TgtPtrAddr,1208 ShadowPtr.TgtPtrContent.data(),1209 ShadowPtr.PtrSize, AsyncInfo);1210 if (Ret != OFFLOAD_SUCCESS) {1211 REPORT("Copying data to device failed.\n");1212 return OFFLOAD_FAIL;1213 }1214 return OFFLOAD_SUCCESS;1215 });1216 if (Ret != OFFLOAD_SUCCESS) {1217 DP("Updating shadow map failed\n");1218 return Ret;1219 }1220 }1221 }1222 1223 if (ArgType & OMP_TGT_MAPTYPE_FROM) {1224 DP("Moving %" PRId64 " bytes (tgt:" DPxMOD ") -> (hst:" DPxMOD ")\n",1225 ArgSize, DPxPTR(TgtPtrBegin), DPxPTR(HstPtrBegin));1226 int Ret = Device.retrieveData(HstPtrBegin, TgtPtrBegin, ArgSize, AsyncInfo,1227 TPR.getEntry());1228 if (Ret != OFFLOAD_SUCCESS) {1229 REPORT("Copying data from device failed.\n");1230 return OFFLOAD_FAIL;1231 }1232 1233 // Wait for device-to-host memcopies for whole struct to complete,1234 // before restoring the correct host pointer/descriptor.1235 if (auto *Entry = TPR.getEntry()) {1236 AsyncInfo.addPostProcessingFunction([=]() -> int {1237 int Ret = Entry->foreachShadowPointerInfo(1238 [&](const ShadowPtrInfoTy &ShadowPtr) {1239 constexpr int64_t VoidPtrSize = sizeof(void *);1240 if (ShadowPtr.PtrSize > VoidPtrSize) {1241 DP("Restoring host descriptor " DPxMOD1242 " to its original content (%" PRId641243 " bytes), containing pointee address " DPxMOD "\n",1244 DPxPTR(ShadowPtr.HstPtrAddr), ShadowPtr.PtrSize,1245 DPxPTR(ShadowPtr.HstPtrContent.data()));1246 } else {1247 DP("Restoring host pointer " DPxMOD1248 " to its original value " DPxMOD "\n",1249 DPxPTR(ShadowPtr.HstPtrAddr),1250 DPxPTR(ShadowPtr.HstPtrContent.data()));1251 }1252 std::memcpy(ShadowPtr.HstPtrAddr, ShadowPtr.HstPtrContent.data(),1253 ShadowPtr.PtrSize);1254 return OFFLOAD_SUCCESS;1255 });1256 Entry->unlock();1257 if (Ret != OFFLOAD_SUCCESS) {1258 DP("Updating shadow map failed\n");1259 return Ret;1260 }1261 return OFFLOAD_SUCCESS;1262 });1263 }1264 }1265 1266 return OFFLOAD_SUCCESS;1267}1268 1269static int targetDataNonContiguous(ident_t *Loc, DeviceTy &Device,1270 void *ArgsBase,1271 __tgt_target_non_contig *NonContig,1272 uint64_t Size, int64_t ArgType,1273 int CurrentDim, int DimSize, uint64_t Offset,1274 AsyncInfoTy &AsyncInfo) {1275 int Ret = OFFLOAD_SUCCESS;1276 if (CurrentDim < DimSize) {1277 for (unsigned int I = 0; I < NonContig[CurrentDim].Count; ++I) {1278 uint64_t CurOffset =1279 (NonContig[CurrentDim].Offset + I) * NonContig[CurrentDim].Stride;1280 // we only need to transfer the first element for the last dimension1281 // since we've already got a contiguous piece.1282 if (CurrentDim != DimSize - 1 || I == 0) {1283 Ret = targetDataNonContiguous(Loc, Device, ArgsBase, NonContig, Size,1284 ArgType, CurrentDim + 1, DimSize,1285 Offset + CurOffset, AsyncInfo);1286 // Stop the whole process if any contiguous piece returns anything1287 // other than OFFLOAD_SUCCESS.1288 if (Ret != OFFLOAD_SUCCESS)1289 return Ret;1290 }1291 }1292 } else {1293 char *Ptr = (char *)ArgsBase + Offset;1294 DP("Transfer of non-contiguous : host ptr " DPxMOD " offset %" PRIu641295 " len %" PRIu64 "\n",1296 DPxPTR(Ptr), Offset, Size);1297 Ret = targetDataContiguous(Loc, Device, ArgsBase, Ptr, Size, ArgType,1298 AsyncInfo);1299 }1300 return Ret;1301}1302 1303static int getNonContigMergedDimension(__tgt_target_non_contig *NonContig,1304 int32_t DimSize) {1305 int RemovedDim = 0;1306 for (int I = DimSize - 1; I > 0; --I) {1307 if (NonContig[I].Count * NonContig[I].Stride == NonContig[I - 1].Stride)1308 RemovedDim++;1309 }1310 return RemovedDim;1311}1312 1313/// Internal function to pass data to/from the target.1314int targetDataUpdate(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,1315 void **ArgsBase, void **Args, int64_t *ArgSizes,1316 int64_t *ArgTypes, map_var_info_t *ArgNames,1317 void **ArgMappers, AsyncInfoTy &AsyncInfo,1318 AttachInfoTy *AttachInfo, bool FromMapper) {1319 // process each input.1320 for (int32_t I = 0; I < ArgNum; ++I) {1321 if ((ArgTypes[I] & OMP_TGT_MAPTYPE_LITERAL) ||1322 (ArgTypes[I] & OMP_TGT_MAPTYPE_PRIVATE))1323 continue;1324 1325 if (ArgMappers && ArgMappers[I]) {1326 // Instead of executing the regular path of targetDataUpdate, call the1327 // targetDataMapper variant which will call targetDataUpdate again1328 // with new arguments.1329 DP("Calling targetDataMapper for the %dth argument\n", I);1330 1331 map_var_info_t ArgName = (!ArgNames) ? nullptr : ArgNames[I];1332 int Ret = targetDataMapper(Loc, Device, ArgsBase[I], Args[I], ArgSizes[I],1333 ArgTypes[I], ArgName, ArgMappers[I], AsyncInfo,1334 targetDataUpdate);1335 1336 if (Ret != OFFLOAD_SUCCESS) {1337 REPORT("Call to targetDataUpdate via targetDataMapper for custom mapper"1338 " failed.\n");1339 return OFFLOAD_FAIL;1340 }1341 1342 // Skip the rest of this function, continue to the next argument.1343 continue;1344 }1345 1346 int Ret = OFFLOAD_SUCCESS;1347 1348 if (ArgTypes[I] & OMP_TGT_MAPTYPE_NON_CONTIG) {1349 __tgt_target_non_contig *NonContig = (__tgt_target_non_contig *)Args[I];1350 int32_t DimSize = ArgSizes[I];1351 uint64_t Size =1352 NonContig[DimSize - 1].Count * NonContig[DimSize - 1].Stride;1353 int32_t MergedDim = getNonContigMergedDimension(NonContig, DimSize);1354 Ret = targetDataNonContiguous(1355 Loc, Device, ArgsBase[I], NonContig, Size, ArgTypes[I],1356 /*current_dim=*/0, DimSize - MergedDim, /*offset=*/0, AsyncInfo);1357 } else {1358 Ret = targetDataContiguous(Loc, Device, ArgsBase[I], Args[I], ArgSizes[I],1359 ArgTypes[I], AsyncInfo);1360 }1361 if (Ret == OFFLOAD_FAIL)1362 return OFFLOAD_FAIL;1363 }1364 return OFFLOAD_SUCCESS;1365}1366 1367static const unsigned LambdaMapping = OMP_TGT_MAPTYPE_PTR_AND_OBJ |1368 OMP_TGT_MAPTYPE_LITERAL |1369 OMP_TGT_MAPTYPE_IMPLICIT;1370static bool isLambdaMapping(int64_t Mapping) {1371 return (Mapping & LambdaMapping) == LambdaMapping;1372}1373 1374namespace {1375/// Find the table information in the map or look it up in the translation1376/// tables.1377TableMap *getTableMap(void *HostPtr) {1378 std::lock_guard<std::mutex> TblMapLock(PM->TblMapMtx);1379 HostPtrToTableMapTy::iterator TableMapIt =1380 PM->HostPtrToTableMap.find(HostPtr);1381 1382 if (TableMapIt != PM->HostPtrToTableMap.end())1383 return &TableMapIt->second;1384 1385 // We don't have a map. So search all the registered libraries.1386 TableMap *TM = nullptr;1387 std::lock_guard<std::mutex> TrlTblLock(PM->TrlTblMtx);1388 for (HostEntriesBeginToTransTableTy::iterator Itr =1389 PM->HostEntriesBeginToTransTable.begin();1390 Itr != PM->HostEntriesBeginToTransTable.end(); ++Itr) {1391 // get the translation table (which contains all the good info).1392 TranslationTable *TransTable = &Itr->second;1393 // iterate over all the host table entries to see if we can locate the1394 // host_ptr.1395 llvm::offloading::EntryTy *Cur = TransTable->HostTable.EntriesBegin;1396 for (uint32_t I = 0; Cur < TransTable->HostTable.EntriesEnd; ++Cur, ++I) {1397 if (Cur->Address != HostPtr)1398 continue;1399 // we got a match, now fill the HostPtrToTableMap so that we1400 // may avoid this search next time.1401 TM = &(PM->HostPtrToTableMap)[HostPtr];1402 TM->Table = TransTable;1403 TM->Index = I;1404 return TM;1405 }1406 }1407 1408 return nullptr;1409}1410 1411/// A class manages private arguments in a target region.1412class PrivateArgumentManagerTy {1413 /// A data structure for the information of first-private arguments. We can1414 /// use this information to optimize data transfer by packing all1415 /// first-private arguments and transfer them all at once.1416 struct FirstPrivateArgInfoTy {1417 /// Host pointer begin1418 char *HstPtrBegin;1419 /// Host pointer end1420 char *HstPtrEnd;1421 /// The index of the element in \p TgtArgs corresponding to the argument1422 int Index;1423 /// Alignment of the entry (base of the entry, not after the entry).1424 uint32_t Alignment;1425 /// Size (without alignment, see padding)1426 uint32_t Size;1427 /// Padding used to align this argument entry, if necessary.1428 uint32_t Padding;1429 /// Host pointer name1430 map_var_info_t HstPtrName = nullptr;1431 /// For corresponding-pointer-initialization: host pointee base address.1432 void *HstPteeBase = nullptr;1433 /// For corresponding-pointer-initialization: host pointee begin address.1434 void *HstPteeBegin = nullptr;1435 /// Whether this argument needs corresponding-pointer-initialization.1436 bool IsCorrespondingPointerInit = false;1437 1438 FirstPrivateArgInfoTy(int Index, void *HstPtr, uint32_t Size,1439 uint32_t Alignment, uint32_t Padding,1440 map_var_info_t HstPtrName = nullptr,1441 void *HstPteeBase = nullptr,1442 void *HstPteeBegin = nullptr,1443 bool IsCorrespondingPointerInit = false)1444 : HstPtrBegin(reinterpret_cast<char *>(HstPtr)),1445 HstPtrEnd(HstPtrBegin + Size), Index(Index), Alignment(Alignment),1446 Size(Size), Padding(Padding), HstPtrName(HstPtrName),1447 HstPteeBase(HstPteeBase), HstPteeBegin(HstPteeBegin),1448 IsCorrespondingPointerInit(IsCorrespondingPointerInit) {}1449 };1450 1451 /// A vector of target pointers for all private arguments1452 SmallVector<void *> TgtPtrs;1453 1454 /// A vector of information of all first-private arguments to be packed1455 SmallVector<FirstPrivateArgInfoTy> FirstPrivateArgInfo;1456 /// Host buffer for all arguments to be packed1457 SmallVector<char> FirstPrivateArgBuffer;1458 /// The total size of all arguments to be packed1459 int64_t FirstPrivateArgSize = 0;1460 1461 /// A reference to the \p DeviceTy object1462 DeviceTy &Device;1463 /// A pointer to a \p AsyncInfoTy object1464 AsyncInfoTy &AsyncInfo;1465 1466 /// \returns the value of the target pointee's base to be used for1467 /// corresponding-pointer-initialization.1468 void *getTargetPointeeBaseForCorrespondingPointerInitialization(1469 void *HstPteeBase, void *HstPteeBegin) {1470 // See if the pointee's begin address has corresponding storage on device.1471 void *TgtPteeBegin = [&]() -> void * {1472 if (!HstPteeBegin) {1473 DP("Corresponding-pointer-initialization: pointee begin address is "1474 "null\n");1475 return nullptr;1476 }1477 1478 return Device.getMappingInfo()1479 .getTgtPtrBegin(HstPteeBegin, /*Size=*/0, /*UpdateRefCount=*/false,1480 /*UseHoldRefCount=*/false)1481 .TargetPointer;1482 }();1483 1484 // If it does, we calculate target pointee base using it, and return it.1485 // Otherwise, we retain the host pointee's base as the target pointee base1486 // of the initialized pointer. It's the user's responsibility to ensure1487 // that if a lookup fails, the host pointee is accessible on the device.1488 return TgtPteeBegin ? calculateTargetPointeeBase(HstPteeBase, HstPteeBegin,1489 TgtPteeBegin)1490 : HstPteeBase;1491 }1492 1493 /// Initialize the source buffer for corresponding-pointer-initialization.1494 ///1495 /// It computes and stores the target pointee base address (or the host1496 /// pointee's base address, if lookup of target pointee fails) to the first1497 /// `sizeof(void*)` bytes of \p Buffer, and for larger pointers1498 /// (Fortran descriptors), the remaining fields of the host descriptor1499 /// \p HstPtr after those `sizeof(void*)` bytes.1500 ///1501 /// Corresponding-pointer-initialization represents the initialization of the1502 /// private version of a base-pointer/referring-pointer on a target construct.1503 ///1504 /// For example, for the following test:1505 /// ```cpp1506 /// int x[10];1507 /// int *px = &x[0];1508 /// ...1509 /// #pragma omp target data map(tofrom:px)1510 /// {1511 /// int **ppx = omp_get_mapped_ptr(&px, omp_get_default_device());1512 /// #pragma omp target map(tofrom:px[1]) is_device_ptr(ppx)1513 /// {1514 /// foo(px, ppx);1515 /// }1516 /// }1517 /// ```1518 /// The following shows a possible way to implement the mapping of `px`,1519 /// which is pre-determined firstprivate and should get initialized1520 /// via corresponding-pointer-initialization:1521 ///1522 /// (A) Possible way to implement the above with PRIVATE | ATTACH:1523 /// ```llvm1524 /// ; maps for px:1525 /// ; &px[0], &px[1], sizeof(px[1]), TO | FROM // (1)1526 /// ; &px, &px[1], sizeof(px), ATTACH // (2)1527 /// ; &px, &px[1], sizeof(px), PRIVATE | ATTACH | PARAM // (3)1528 /// call... @__omp_outlined...(ptr %px, ptr %ppx)1529 /// define ... @__omp_outlined(ptr %px, ptr %ppx) {...1530 /// foo(%px, %ppx)1531 /// ...}1532 /// ```1533 /// `(1)` maps the pointee `px[1].1534 /// `(2)` attaches it to the mapped version of `px`. It can be controlled by1535 /// the user based on the `attach(auto/always/never)` map-type modifier.1536 /// `(3)` privatizes and initializes the private pointer `px`, and passes it1537 /// into the kernel as the argument `%px`. Can be skipped if `px` is not1538 /// referenced in the target construct.1539 ///1540 /// While this method is not too beneficial compared to just doing the1541 /// initialization in the body of the kernel, like:1542 /// (B) Possible way to implement the above without PRIVATE | ATTACH:1543 /// ```llvm1544 /// ; maps for px:1545 /// ; &px[0], &px[1], sizeof(px[1]), TO | FROM | PARAM // (4)1546 /// ; &px, &px[1], sizeof(px), ATTACH // (5)1547 /// call... @__omp_outlined...(ptr %px0, ptr %ppx)1548 /// define ... __omp_outlined...(ptr %px0, ptr %ppx) {1549 /// %px = alloca ptr;1550 /// store ptr %px0, ptr %px1551 /// foo(%px, %ppx)1552 /// }1553 /// ```1554 ///1555 /// (B) is not so convenient for Fortran descriptors, because in1556 /// addition to the lookup, the remaining fields of the descriptor have1557 /// to be passed into the kernel to initialize the private copy, which1558 /// makes (A) a cleaner option for them. e.g.1559 /// ```f901560 /// integer, pointer :: p(:)1561 /// !$omp target map(p(1))1562 /// ```1563 ///1564 /// (C) Possible mapping for the above Fortran test using PRIVATE | ATTACH:1565 /// ```llvm1566 /// ; maps for p:1567 /// ; &p(1), &p(1), sizeof(p(1)), TO | FROM1568 /// ; &ref_ptr(p), &p(1), sizeof(ref_ptr(p)), ATTACH1569 /// ; &ref_ptr(p), &p(1), sizeof(ref_ptr(p)), PRIVATE | ATTACH | PARAM1570 /// call... @__omp_outlined...(ptr %ref_ptr_of_p)1571 void initBufferForCorrespondingPointerInitialization(char *Buffer,1572 void *HstPtr,1573 int64_t HstPtrSize,1574 void *HstPteeBase,1575 void *HstPteeBegin) {1576 constexpr int64_t VoidPtrSize = sizeof(void *);1577 assert(HstPtrSize >= VoidPtrSize &&1578 "corresponding-pointer-initialization: pointer size is too small");1579 1580 void *TgtPteeBase =1581 getTargetPointeeBaseForCorrespondingPointerInitialization(HstPteeBase,1582 HstPteeBegin);1583 1584 // Store the target pointee base address to the first VoidPtrSize bytes1585 DP("Initializing corresponding-pointer-initialization source buffer "1586 "for " DPxMOD ", with pointee base " DPxMOD "\n",1587 DPxPTR(HstPtr), DPxPTR(TgtPteeBase));1588 std::memcpy(Buffer, &TgtPteeBase, VoidPtrSize);1589 if (HstPtrSize <= VoidPtrSize)1590 return;1591 1592 // For Fortran descriptors, copy the remaining descriptor fields from host1593 uint64_t HstDescriptorFieldsSize = HstPtrSize - VoidPtrSize;1594 void *HstDescriptorFieldsAddr = static_cast<char *>(HstPtr) + VoidPtrSize;1595 DP("Copying %" PRId641596 " bytes of descriptor fields into corresponding-pointer-initialization "1597 "buffer at offset %" PRId64 ", from " DPxMOD "\n",1598 HstDescriptorFieldsSize, VoidPtrSize, DPxPTR(HstDescriptorFieldsAddr));1599 std::memcpy(Buffer + VoidPtrSize, HstDescriptorFieldsAddr,1600 HstDescriptorFieldsSize);1601 }1602 1603 /// Helper function to create and initialize a buffer to be used as the source1604 /// for corresponding-pointer-initialization.1605 void *createAndInitSourceBufferForCorrespondingPointerInitialization(1606 void *HstPtr, int64_t HstPtrSize, void *HstPteeBase, void *HstPteeBegin) {1607 char *Buffer = getOrCreateSourceBufferForSubmitData(AsyncInfo, HstPtrSize);1608 initBufferForCorrespondingPointerInitialization(Buffer, HstPtr, HstPtrSize,1609 HstPteeBase, HstPteeBegin);1610 return Buffer;1611 }1612 1613 // TODO: What would be the best value here? Should we make it configurable?1614 // If the size is larger than this threshold, we will allocate and transfer it1615 // immediately instead of packing it.1616 static constexpr const int64_t FirstPrivateArgSizeThreshold = 1024;1617 1618public:1619 /// Constructor1620 PrivateArgumentManagerTy(DeviceTy &Dev, AsyncInfoTy &AsyncInfo)1621 : Device(Dev), AsyncInfo(AsyncInfo) {}1622 1623 /// Add a private argument1624 int addArg(void *HstPtr, int64_t ArgSize, int64_t ArgOffset,1625 bool IsFirstPrivate, void *&TgtPtr, int TgtArgsIndex,1626 map_var_info_t HstPtrName = nullptr,1627 const bool AllocImmediately = false, void *HstPteeBase = nullptr,1628 void *HstPteeBegin = nullptr,1629 bool IsCorrespondingPointerInit = false) {1630 // If the argument is not first-private, or its size is greater than a1631 // predefined threshold, we will allocate memory and issue the transfer1632 // immediately.1633 if (ArgSize > FirstPrivateArgSizeThreshold || !IsFirstPrivate ||1634 AllocImmediately) {1635 TgtPtr = Device.allocData(ArgSize, HstPtr);1636 if (!TgtPtr) {1637 DP("Data allocation for %sprivate array " DPxMOD " failed.\n",1638 (IsFirstPrivate ? "first-" : ""), DPxPTR(HstPtr));1639 return OFFLOAD_FAIL;1640 }1641#ifdef OMPTARGET_DEBUG1642 void *TgtPtrBase = (void *)((intptr_t)TgtPtr + ArgOffset);1643 DP("Allocated %" PRId64 " bytes of target memory at " DPxMOD1644 " for %sprivate array " DPxMOD " - pushing target argument " DPxMOD1645 "\n",1646 ArgSize, DPxPTR(TgtPtr), (IsFirstPrivate ? "first-" : ""),1647 DPxPTR(HstPtr), DPxPTR(TgtPtrBase));1648#endif1649 // If first-private, copy data from host1650 if (IsFirstPrivate) {1651 DP("Submitting firstprivate data to the device.\n");1652 1653 // The source value used for corresponding-pointer-initialization1654 // is different vs regular firstprivates.1655 void *DataSource =1656 IsCorrespondingPointerInit1657 ? createAndInitSourceBufferForCorrespondingPointerInitialization(1658 HstPtr, ArgSize, HstPteeBase, HstPteeBegin)1659 : HstPtr;1660 int Ret = Device.submitData(TgtPtr, DataSource, ArgSize, AsyncInfo);1661 if (Ret != OFFLOAD_SUCCESS) {1662 DP("Copying %s data to device failed.\n",1663 IsCorrespondingPointerInit ? "corresponding-pointer-initialization"1664 : "firstprivate");1665 return OFFLOAD_FAIL;1666 }1667 }1668 TgtPtrs.push_back(TgtPtr);1669 } else {1670 DP("Firstprivate array " DPxMOD " of size %" PRId64 " will be packed\n",1671 DPxPTR(HstPtr), ArgSize);1672 // When reach this point, the argument must meet all following1673 // requirements:1674 // 1. Its size does not exceed the threshold (see the comment for1675 // FirstPrivateArgSizeThreshold);1676 // 2. It must be first-private (needs to be mapped to target device).1677 // We will pack all this kind of arguments to transfer them all at once1678 // to reduce the number of data transfer. We will not take1679 // non-first-private arguments, aka. private arguments that doesn't need1680 // to be mapped to target device, into account because data allocation1681 // can be very efficient with memory manager.1682 1683 // Placeholder value1684 TgtPtr = nullptr;1685 auto *LastFPArgInfo =1686 FirstPrivateArgInfo.empty() ? nullptr : &FirstPrivateArgInfo.back();1687 1688 // Compute the start alignment of this entry, add padding if necessary.1689 // TODO: Consider sorting instead.1690 uint32_t Padding = 0;1691 uint32_t StartAlignment =1692 LastFPArgInfo ? LastFPArgInfo->Alignment : MaxAlignment;1693 if (LastFPArgInfo) {1694 // Check if we keep the start alignment or if it is shrunk due to the1695 // size of the last element.1696 uint32_t Offset = LastFPArgInfo->Size % StartAlignment;1697 if (Offset)1698 StartAlignment = Offset;1699 // We only need as much alignment as the host pointer had (since we1700 // don't know the alignment information from the source we might end up1701 // overaligning accesses but not too much).1702 uint32_t RequiredAlignment =1703 llvm::bit_floor(getPartialStructRequiredAlignment(HstPtr));1704 if (RequiredAlignment > StartAlignment) {1705 Padding = RequiredAlignment - StartAlignment;1706 StartAlignment = RequiredAlignment;1707 }1708 }1709 1710 FirstPrivateArgInfo.emplace_back(1711 TgtArgsIndex, HstPtr, ArgSize, StartAlignment, Padding, HstPtrName,1712 HstPteeBase, HstPteeBegin, IsCorrespondingPointerInit);1713 1714 FirstPrivateArgSize += Padding + ArgSize;1715 }1716 1717 return OFFLOAD_SUCCESS;1718 }1719 1720 /// Pack first-private arguments, replace place holder pointers in \p TgtArgs,1721 /// and start the transfer.1722 int packAndTransfer(SmallVector<void *> &TgtArgs) {1723 if (!FirstPrivateArgInfo.empty()) {1724 assert(FirstPrivateArgSize != 0 &&1725 "FirstPrivateArgSize is 0 but FirstPrivateArgInfo is empty");1726 FirstPrivateArgBuffer.resize(FirstPrivateArgSize, 0);1727 auto *Itr = FirstPrivateArgBuffer.begin();1728 // Copy all host data to this buffer1729 for (FirstPrivateArgInfoTy &Info : FirstPrivateArgInfo) {1730 // First pad the pointer as we (have to) pad it on the device too.1731 Itr = std::next(Itr, Info.Padding);1732 1733 if (Info.IsCorrespondingPointerInit)1734 initBufferForCorrespondingPointerInitialization(1735 &*Itr, Info.HstPtrBegin, Info.Size, Info.HstPteeBase,1736 Info.HstPteeBegin);1737 else1738 std::copy(Info.HstPtrBegin, Info.HstPtrEnd, Itr);1739 Itr = std::next(Itr, Info.Size);1740 }1741 // Allocate target memory1742 void *TgtPtr =1743 Device.allocData(FirstPrivateArgSize, FirstPrivateArgBuffer.data());1744 if (TgtPtr == nullptr) {1745 DP("Failed to allocate target memory for private arguments.\n");1746 return OFFLOAD_FAIL;1747 }1748 TgtPtrs.push_back(TgtPtr);1749 DP("Allocated %" PRId64 " bytes of target memory at " DPxMOD "\n",1750 FirstPrivateArgSize, DPxPTR(TgtPtr));1751 // Transfer data to target device1752 int Ret = Device.submitData(TgtPtr, FirstPrivateArgBuffer.data(),1753 FirstPrivateArgSize, AsyncInfo);1754 if (Ret != OFFLOAD_SUCCESS) {1755 DP("Failed to submit data of private arguments.\n");1756 return OFFLOAD_FAIL;1757 }1758 // Fill in all placeholder pointers1759 auto TP = reinterpret_cast<uintptr_t>(TgtPtr);1760 for (FirstPrivateArgInfoTy &Info : FirstPrivateArgInfo) {1761 void *&Ptr = TgtArgs[Info.Index];1762 assert(Ptr == nullptr && "Target pointer is already set by mistaken");1763 // Pad the device pointer to get the right alignment.1764 TP += Info.Padding;1765 Ptr = reinterpret_cast<void *>(TP);1766 TP += Info.Size;1767 DP("Firstprivate array " DPxMOD " of size %" PRId64 " mapped to " DPxMOD1768 "\n",1769 DPxPTR(Info.HstPtrBegin), Info.HstPtrEnd - Info.HstPtrBegin,1770 DPxPTR(Ptr));1771 }1772 }1773 1774 return OFFLOAD_SUCCESS;1775 }1776 1777 /// Free all target memory allocated for private arguments1778 int free() {1779 for (void *P : TgtPtrs) {1780 int Ret = Device.deleteData(P);1781 if (Ret != OFFLOAD_SUCCESS) {1782 DP("Deallocation of (first-)private arrays failed.\n");1783 return OFFLOAD_FAIL;1784 }1785 }1786 1787 TgtPtrs.clear();1788 1789 return OFFLOAD_SUCCESS;1790 }1791};1792 1793/// Process data before launching the kernel, including calling targetDataBegin1794/// to map and transfer data to target device, transferring (first-)private1795/// variables.1796static int processDataBefore(ident_t *Loc, int64_t DeviceId, void *HostPtr,1797 int32_t ArgNum, void **ArgBases, void **Args,1798 int64_t *ArgSizes, int64_t *ArgTypes,1799 map_var_info_t *ArgNames, void **ArgMappers,1800 SmallVector<void *> &TgtArgs,1801 SmallVector<ptrdiff_t> &TgtOffsets,1802 PrivateArgumentManagerTy &PrivateArgumentManager,1803 AsyncInfoTy &AsyncInfo) {1804 1805 auto DeviceOrErr = PM->getDevice(DeviceId);1806 if (!DeviceOrErr)1807 FATAL_MESSAGE(DeviceId, "%s", toString(DeviceOrErr.takeError()).c_str());1808 1809 // Create AttachInfo for tracking any ATTACH entries, or new-allocations1810 // when handling the "begin" mapping for a target constructs.1811 AttachInfoTy AttachInfo;1812 1813 int Ret = targetDataBegin(Loc, *DeviceOrErr, ArgNum, ArgBases, Args, ArgSizes,1814 ArgTypes, ArgNames, ArgMappers, AsyncInfo,1815 &AttachInfo, false /*FromMapper=*/);1816 if (Ret != OFFLOAD_SUCCESS) {1817 REPORT("Call to targetDataBegin failed, abort target.\n");1818 return OFFLOAD_FAIL;1819 }1820 1821 // Process collected ATTACH entries1822 if (!AttachInfo.AttachEntries.empty()) {1823 Ret = processAttachEntries(*DeviceOrErr, AttachInfo, AsyncInfo);1824 if (Ret != OFFLOAD_SUCCESS) {1825 REPORT("Failed to process ATTACH entries.\n");1826 return OFFLOAD_FAIL;1827 }1828 }1829 1830 // List of (first-)private arrays allocated for this target region1831 SmallVector<int> TgtArgsPositions(ArgNum, -1);1832 1833 for (int32_t I = 0; I < ArgNum; ++I) {1834 if (!(ArgTypes[I] & OMP_TGT_MAPTYPE_TARGET_PARAM)) {1835 // This is not a target parameter, do not push it into TgtArgs.1836 // Check for lambda mapping.1837 if (isLambdaMapping(ArgTypes[I])) {1838 assert((ArgTypes[I] & OMP_TGT_MAPTYPE_MEMBER_OF) &&1839 "PTR_AND_OBJ must be also MEMBER_OF.");1840 unsigned Idx = getParentIndex(ArgTypes[I]);1841 int TgtIdx = TgtArgsPositions[Idx];1842 assert(TgtIdx != -1 && "Base address must be translated already.");1843 // The parent lambda must be processed already and it must be the last1844 // in TgtArgs and TgtOffsets arrays.1845 void *HstPtrVal = Args[I];1846 void *HstPtrBegin = ArgBases[I];1847 void *HstPtrBase = Args[Idx];1848 void *TgtPtrBase =1849 (void *)((intptr_t)TgtArgs[TgtIdx] + TgtOffsets[TgtIdx]);1850 DP("Parent lambda base " DPxMOD "\n", DPxPTR(TgtPtrBase));1851 uint64_t Delta = (uint64_t)HstPtrBegin - (uint64_t)HstPtrBase;1852 void *TgtPtrBegin = (void *)((uintptr_t)TgtPtrBase + Delta);1853 void *&PointerTgtPtrBegin = AsyncInfo.getVoidPtrLocation();1854 TargetPointerResultTy TPR =1855 DeviceOrErr->getMappingInfo().getTgtPtrBegin(1856 HstPtrVal, ArgSizes[I], /*UpdateRefCount=*/false,1857 /*UseHoldRefCount=*/false);1858 PointerTgtPtrBegin = TPR.TargetPointer;1859 if (!TPR.isPresent()) {1860 DP("No lambda captured variable mapped (" DPxMOD ") - ignored\n",1861 DPxPTR(HstPtrVal));1862 continue;1863 }1864 if (TPR.Flags.IsHostPointer) {1865 DP("Unified memory is active, no need to map lambda captured"1866 "variable (" DPxMOD ")\n",1867 DPxPTR(HstPtrVal));1868 continue;1869 }1870 DP("Update lambda reference (" DPxMOD ") -> [" DPxMOD "]\n",1871 DPxPTR(PointerTgtPtrBegin), DPxPTR(TgtPtrBegin));1872 Ret =1873 DeviceOrErr->submitData(TgtPtrBegin, &PointerTgtPtrBegin,1874 sizeof(void *), AsyncInfo, TPR.getEntry());1875 if (Ret != OFFLOAD_SUCCESS) {1876 REPORT("Copying data to device failed.\n");1877 return OFFLOAD_FAIL;1878 }1879 }1880 continue;1881 }1882 void *HstPtrBegin = Args[I];1883 void *HstPtrBase = ArgBases[I];1884 void *TgtPtrBegin;1885 map_var_info_t HstPtrName = (!ArgNames) ? nullptr : ArgNames[I];1886 ptrdiff_t TgtBaseOffset;1887 TargetPointerResultTy TPR;1888 if (ArgTypes[I] & OMP_TGT_MAPTYPE_LITERAL) {1889 DP("Forwarding first-private value " DPxMOD " to the target construct\n",1890 DPxPTR(HstPtrBase));1891 TgtPtrBegin = HstPtrBase;1892 TgtBaseOffset = 0;1893 } else if (ArgTypes[I] & OMP_TGT_MAPTYPE_PRIVATE) {1894 // For cases like:1895 // ```1896 // int *p = ...;1897 // #pragma omp target map(p[0:10])1898 // ```1899 // `p` is predetermined firstprivate on the target construct, and the1900 // method to determine the initial value of the private copy on the1901 // device is called "corresponding-pointer-initialization".1902 //1903 // Such firstprivate pointers that need1904 // corresponding-pointer-initialization are represented using the1905 // `PRIVATE | ATTACH` map-types, in contrast to regular firstprivate1906 // entries, which use `PRIVATE | TO`. The structure of these1907 // `PRIVATE | ATTACH` entries is the same as the non-private1908 // `ATTACH` entries used to represent pointer-attachments, i.e.:1909 // ```1910 // &hst_ptr_base/begin, &hst_ptee_begin, sizeof(hst_ptr)1911 // ```1912 const bool IsAttach = (ArgTypes[I] & OMP_TGT_MAPTYPE_ATTACH);1913 void *HstPteeBase = nullptr;1914 void *HstPteeBegin = nullptr;1915 if (IsAttach) {1916 // For corresponding-pointer-initialization, Args[I] is HstPteeBegin,1917 // and ArgBases[I] is both HstPtrBase/HstPtrBegin.1918 HstPteeBase = *reinterpret_cast<void **>(HstPtrBase);1919 HstPteeBegin = Args[I];1920 HstPtrBegin = ArgBases[I];1921 }1922 TgtBaseOffset = (intptr_t)HstPtrBase - (intptr_t)HstPtrBegin;1923 // Corresponding-pointer-initialization is a special case of firstprivate,1924 // since it also involves initializing the private pointer.1925 const bool IsFirstPrivate =1926 (ArgTypes[I] & OMP_TGT_MAPTYPE_TO) || IsAttach;1927 1928 // If there is a next argument and it depends on the current one, we need1929 // to allocate the private memory immediately. If this is not the case,1930 // then the argument can be marked for optimization and packed with the1931 // other privates.1932 const bool AllocImmediately =1933 (I < ArgNum - 1 && (ArgTypes[I + 1] & OMP_TGT_MAPTYPE_MEMBER_OF));1934 Ret = PrivateArgumentManager.addArg(1935 HstPtrBegin, ArgSizes[I], TgtBaseOffset, IsFirstPrivate, TgtPtrBegin,1936 /*TgtArgsIndex=*/TgtArgs.size(), HstPtrName, AllocImmediately,1937 HstPteeBase, HstPteeBegin, /*IsCorrespondingPointerInit=*/IsAttach);1938 if (Ret != OFFLOAD_SUCCESS) {1939 REPORT("Failed to process %s%sprivate argument " DPxMOD "\n",1940 IsAttach ? "corresponding-pointer-initialization " : "",1941 (IsFirstPrivate ? "first-" : ""), DPxPTR(HstPtrBegin));1942 return OFFLOAD_FAIL;1943 }1944 } else {1945 if (ArgTypes[I] & OMP_TGT_MAPTYPE_PTR_AND_OBJ)1946 HstPtrBase = *reinterpret_cast<void **>(HstPtrBase);1947 TPR = DeviceOrErr->getMappingInfo().getTgtPtrBegin(1948 HstPtrBegin, ArgSizes[I],1949 /*UpdateRefCount=*/false,1950 /*UseHoldRefCount=*/false);1951 TgtPtrBegin = TPR.TargetPointer;1952 TgtBaseOffset = (intptr_t)HstPtrBase - (intptr_t)HstPtrBegin;1953#ifdef OMPTARGET_DEBUG1954 void *TgtPtrBase = (void *)((intptr_t)TgtPtrBegin + TgtBaseOffset);1955 DP("Obtained target argument " DPxMOD " from host pointer " DPxMOD "\n",1956 DPxPTR(TgtPtrBase), DPxPTR(HstPtrBegin));1957#endif1958 }1959 TgtArgsPositions[I] = TgtArgs.size();1960 TgtArgs.push_back(TgtPtrBegin);1961 TgtOffsets.push_back(TgtBaseOffset);1962 }1963 1964 assert(TgtArgs.size() == TgtOffsets.size() &&1965 "Size mismatch in arguments and offsets");1966 1967 // Pack and transfer first-private arguments1968 Ret = PrivateArgumentManager.packAndTransfer(TgtArgs);1969 if (Ret != OFFLOAD_SUCCESS) {1970 DP("Failed to pack and transfer first private arguments\n");1971 return OFFLOAD_FAIL;1972 }1973 1974 return OFFLOAD_SUCCESS;1975}1976 1977/// Process data after launching the kernel, including transferring data back to1978/// host if needed and deallocating target memory of (first-)private variables.1979static int processDataAfter(ident_t *Loc, int64_t DeviceId, void *HostPtr,1980 int32_t ArgNum, void **ArgBases, void **Args,1981 int64_t *ArgSizes, int64_t *ArgTypes,1982 map_var_info_t *ArgNames, void **ArgMappers,1983 PrivateArgumentManagerTy &PrivateArgumentManager,1984 AsyncInfoTy &AsyncInfo) {1985 1986 auto DeviceOrErr = PM->getDevice(DeviceId);1987 if (!DeviceOrErr)1988 FATAL_MESSAGE(DeviceId, "%s", toString(DeviceOrErr.takeError()).c_str());1989 1990 // Move data from device.1991 int Ret = targetDataEnd(Loc, *DeviceOrErr, ArgNum, ArgBases, Args, ArgSizes,1992 ArgTypes, ArgNames, ArgMappers, AsyncInfo);1993 if (Ret != OFFLOAD_SUCCESS) {1994 REPORT("Call to targetDataEnd failed, abort target.\n");1995 return OFFLOAD_FAIL;1996 }1997 1998 // Free target memory for private arguments after synchronization.1999 // TODO: We might want to remove `mutable` in the future by not changing the2000 // captured variables somehow.2001 AsyncInfo.addPostProcessingFunction(2002 [PrivateArgumentManager =2003 std::move(PrivateArgumentManager)]() mutable -> int {2004 int Ret = PrivateArgumentManager.free();2005 if (Ret != OFFLOAD_SUCCESS) {2006 REPORT("Failed to deallocate target memory for private args\n");2007 return OFFLOAD_FAIL;2008 }2009 return Ret;2010 });2011 2012 return OFFLOAD_SUCCESS;2013}2014} // namespace2015 2016/// performs the same actions as data_begin in case arg_num is2017/// non-zero and initiates run of the offloaded region on the target platform;2018/// if arg_num is non-zero after the region execution is done it also2019/// performs the same action as data_update and data_end above. This function2020/// returns 0 if it was able to transfer the execution to a target and an2021/// integer different from zero otherwise.2022int target(ident_t *Loc, DeviceTy &Device, void *HostPtr,2023 KernelArgsTy &KernelArgs, AsyncInfoTy &AsyncInfo) {2024 int32_t DeviceId = Device.DeviceID;2025 TableMap *TM = getTableMap(HostPtr);2026 // No map for this host pointer found!2027 if (!TM) {2028 REPORT("Host ptr " DPxMOD " does not have a matching target pointer.\n",2029 DPxPTR(HostPtr));2030 return OFFLOAD_FAIL;2031 }2032 2033 // get target table.2034 __tgt_target_table *TargetTable = nullptr;2035 {2036 std::lock_guard<std::mutex> TrlTblLock(PM->TrlTblMtx);2037 assert(TM->Table->TargetsTable.size() > (size_t)DeviceId &&2038 "Not expecting a device ID outside the table's bounds!");2039 TargetTable = TM->Table->TargetsTable[DeviceId];2040 }2041 assert(TargetTable && "Global data has not been mapped\n");2042 2043 DP("loop trip count is %" PRIu64 ".\n", KernelArgs.Tripcount);2044 2045 // We need to keep bases and offsets separate. Sometimes (e.g. in OpenCL) we2046 // need to manifest base pointers prior to launching a kernel. Even if we have2047 // mapped an object only partially, e.g. A[N:M], although the kernel is2048 // expected to access elements starting at address &A[N] and beyond, we still2049 // need to manifest the base of the array &A[0]. In other cases, e.g. the COI2050 // API, we need the begin address itself, i.e. &A[N], as the API operates on2051 // begin addresses, not bases. That's why we pass args and offsets as two2052 // separate entities so that each plugin can do what it needs. This behavior2053 // was introduced via https://reviews.llvm.org/D33028 and commit 1546d319244c.2054 SmallVector<void *> TgtArgs;2055 SmallVector<ptrdiff_t> TgtOffsets;2056 2057 PrivateArgumentManagerTy PrivateArgumentManager(Device, AsyncInfo);2058 2059 int NumClangLaunchArgs = KernelArgs.NumArgs;2060 int Ret = OFFLOAD_SUCCESS;2061 if (NumClangLaunchArgs) {2062 // Process data, such as data mapping, before launching the kernel2063 Ret = processDataBefore(Loc, DeviceId, HostPtr, NumClangLaunchArgs,2064 KernelArgs.ArgBasePtrs, KernelArgs.ArgPtrs,2065 KernelArgs.ArgSizes, KernelArgs.ArgTypes,2066 KernelArgs.ArgNames, KernelArgs.ArgMappers, TgtArgs,2067 TgtOffsets, PrivateArgumentManager, AsyncInfo);2068 if (Ret != OFFLOAD_SUCCESS) {2069 REPORT("Failed to process data before launching the kernel.\n");2070 return OFFLOAD_FAIL;2071 }2072 2073 // Clang might pass more values via the ArgPtrs to the runtime that we pass2074 // on to the kernel.2075 // TODO: Next time we adjust the KernelArgsTy we should introduce a new2076 // NumKernelArgs field.2077 KernelArgs.NumArgs = TgtArgs.size();2078 }2079 2080 // Launch device execution.2081 void *TgtEntryPtr = TargetTable->EntriesBegin[TM->Index].Address;2082 DP("Launching target execution %s with pointer " DPxMOD " (index=%d).\n",2083 TargetTable->EntriesBegin[TM->Index].SymbolName, DPxPTR(TgtEntryPtr),2084 TM->Index);2085 2086 {2087 assert(KernelArgs.NumArgs == TgtArgs.size() && "Argument count mismatch!");2088 TIMESCOPE_WITH_DETAILS_AND_IDENT(2089 "Kernel Target",2090 "NumArguments=" + std::to_string(KernelArgs.NumArgs) +2091 ";NumTeams=" + std::to_string(KernelArgs.NumTeams[0]) +2092 ";TripCount=" + std::to_string(KernelArgs.Tripcount),2093 Loc);2094 2095#ifdef OMPT_SUPPORT2096 /// RAII to establish tool anchors before and after kernel launch2097 int32_t NumTeams = KernelArgs.NumTeams[0];2098 // No need to guard this with OMPT_IF_BUILT2099 InterfaceRAII TargetSubmitRAII(2100 RegionInterface.getCallbacks<ompt_callback_target_submit>(), NumTeams);2101#endif2102 2103 Ret = Device.launchKernel(TgtEntryPtr, TgtArgs.data(), TgtOffsets.data(),2104 KernelArgs, AsyncInfo);2105 }2106 2107 if (Ret != OFFLOAD_SUCCESS) {2108 REPORT("Executing target region abort target.\n");2109 return OFFLOAD_FAIL;2110 }2111 2112 if (NumClangLaunchArgs) {2113 // Transfer data back and deallocate target memory for (first-)private2114 // variables2115 Ret = processDataAfter(Loc, DeviceId, HostPtr, NumClangLaunchArgs,2116 KernelArgs.ArgBasePtrs, KernelArgs.ArgPtrs,2117 KernelArgs.ArgSizes, KernelArgs.ArgTypes,2118 KernelArgs.ArgNames, KernelArgs.ArgMappers,2119 PrivateArgumentManager, AsyncInfo);2120 if (Ret != OFFLOAD_SUCCESS) {2121 REPORT("Failed to process data after launching the kernel.\n");2122 return OFFLOAD_FAIL;2123 }2124 }2125 2126 return OFFLOAD_SUCCESS;2127}2128 2129/// Enables the record replay mechanism by pre-allocating MemorySize2130/// and informing the record-replayer of whether to store the output2131/// in some file.2132int target_activate_rr(DeviceTy &Device, uint64_t MemorySize, void *VAddr,2133 bool IsRecord, bool SaveOutput,2134 uint64_t &ReqPtrArgOffset) {2135 return Device.RTL->initialize_record_replay(Device.DeviceID, MemorySize,2136 VAddr, IsRecord, SaveOutput,2137 ReqPtrArgOffset);2138}2139 2140/// Executes a kernel using pre-recorded information for loading to2141/// device memory to launch the target kernel with the pre-recorded2142/// configuration.2143int target_replay(ident_t *Loc, DeviceTy &Device, void *HostPtr,2144 void *DeviceMemory, int64_t DeviceMemorySize, void **TgtArgs,2145 ptrdiff_t *TgtOffsets, int32_t NumArgs, int32_t NumTeams,2146 int32_t ThreadLimit, uint64_t LoopTripCount,2147 AsyncInfoTy &AsyncInfo) {2148 int32_t DeviceId = Device.DeviceID;2149 TableMap *TM = getTableMap(HostPtr);2150 // Fail if the table map fails to find the target kernel pointer for the2151 // provided host pointer.2152 if (!TM) {2153 REPORT("Host ptr " DPxMOD " does not have a matching target pointer.\n",2154 DPxPTR(HostPtr));2155 return OFFLOAD_FAIL;2156 }2157 2158 // Retrieve the target table of offloading entries.2159 __tgt_target_table *TargetTable = nullptr;2160 {2161 std::lock_guard<std::mutex> TrlTblLock(PM->TrlTblMtx);2162 assert(TM->Table->TargetsTable.size() > (size_t)DeviceId &&2163 "Not expecting a device ID outside the table's bounds!");2164 TargetTable = TM->Table->TargetsTable[DeviceId];2165 }2166 assert(TargetTable && "Global data has not been mapped\n");2167 2168 // Retrieve the target kernel pointer, allocate and store the recorded device2169 // memory data, and launch device execution.2170 void *TgtEntryPtr = TargetTable->EntriesBegin[TM->Index].Address;2171 DP("Launching target execution %s with pointer " DPxMOD " (index=%d).\n",2172 TargetTable->EntriesBegin[TM->Index].SymbolName, DPxPTR(TgtEntryPtr),2173 TM->Index);2174 2175 void *TgtPtr = Device.allocData(DeviceMemorySize, /*HstPtr=*/nullptr,2176 TARGET_ALLOC_DEFAULT);2177 Device.submitData(TgtPtr, DeviceMemory, DeviceMemorySize, AsyncInfo);2178 2179 KernelArgsTy KernelArgs{};2180 KernelArgs.Version = OMP_KERNEL_ARG_VERSION;2181 KernelArgs.NumArgs = NumArgs;2182 KernelArgs.Tripcount = LoopTripCount;2183 KernelArgs.NumTeams[0] = NumTeams;2184 KernelArgs.ThreadLimit[0] = ThreadLimit;2185 2186 int Ret = Device.launchKernel(TgtEntryPtr, TgtArgs, TgtOffsets, KernelArgs,2187 AsyncInfo);2188 2189 if (Ret != OFFLOAD_SUCCESS) {2190 REPORT("Executing target region abort target.\n");2191 return OFFLOAD_FAIL;2192 }2193 2194 return OFFLOAD_SUCCESS;2195}2196