2216 lines · cpp
1//===- PluginInterface.cpp - Target independent plugin device interface ---===//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//===----------------------------------------------------------------------===//10 11#include "PluginInterface.h"12 13#include "Shared/APITypes.h"14#include "Shared/Debug.h"15#include "Shared/Environment.h"16 17#include "ErrorReporting.h"18#include "GlobalHandler.h"19#include "JIT.h"20#include "Shared/Utils.h"21#include "Utils/ELF.h"22#include "omptarget.h"23 24#ifdef OMPT_SUPPORT25#include "OpenMP/OMPT/Callback.h"26#include "omp-tools.h"27#endif28 29#include "llvm/Bitcode/BitcodeReader.h"30#include "llvm/Frontend/OpenMP/OMPConstants.h"31#include "llvm/Support/Error.h"32#include "llvm/Support/JSON.h"33#include "llvm/Support/MathExtras.h"34#include "llvm/Support/MemoryBuffer.h"35#include "llvm/Support/Signals.h"36#include "llvm/Support/raw_ostream.h"37 38#include <cstdint>39#include <limits>40 41using namespace llvm;42using namespace omp;43using namespace target;44using namespace plugin;45using namespace error;46 47// TODO: Fix any thread safety issues for multi-threaded kernel recording.48namespace llvm::omp::target::plugin {49struct RecordReplayTy {50 51 // Describes the state of the record replay mechanism.52 enum RRStatusTy { RRDeactivated = 0, RRRecording, RRReplaying };53 54private:55 // Memory pointers for recording, replaying memory.56 void *MemoryStart = nullptr;57 void *MemoryPtr = nullptr;58 size_t MemorySize = 0;59 size_t TotalSize = 0;60 GenericDeviceTy *Device = nullptr;61 std::mutex AllocationLock;62 63 RRStatusTy Status = RRDeactivated;64 bool ReplaySaveOutput = false;65 bool UsedVAMap = false;66 uintptr_t MemoryOffset = 0;67 68 // A list of all globals mapped to the device.69 struct GlobalEntry {70 const char *Name;71 uint64_t Size;72 void *Addr;73 };74 llvm::SmallVector<GlobalEntry> GlobalEntries{};75 76 Expected<void *> suggestAddress(uint64_t MaxMemoryAllocation) {77 // Get a valid pointer address for this system78 auto AddrOrErr =79 Device->allocate(1024, /*HstPtr=*/nullptr, TARGET_ALLOC_DEFAULT);80 if (!AddrOrErr)81 return AddrOrErr.takeError();82 83 void *Addr = *AddrOrErr;84 if (auto Err = Device->free(Addr))85 return std::move(Err);86 87 // Align Address to MaxMemoryAllocation88 Addr = (void *)utils::alignPtr((Addr), MaxMemoryAllocation);89 return Addr;90 }91 92 Error preAllocateVAMemory(uint64_t MaxMemoryAllocation, void *VAddr) {93 size_t ASize = MaxMemoryAllocation;94 95 if (!VAddr && isRecording()) {96 auto VAddrOrErr = suggestAddress(MaxMemoryAllocation);97 if (!VAddrOrErr)98 return VAddrOrErr.takeError();99 VAddr = *VAddrOrErr;100 }101 102 DP("Request %ld bytes allocated at %p\n", MaxMemoryAllocation, VAddr);103 104 if (auto Err = Device->memoryVAMap(&MemoryStart, VAddr, &ASize))105 return Err;106 107 if (isReplaying() && VAddr != MemoryStart) {108 return Plugin::error(ErrorCode::INVALID_ARGUMENT,109 "record-Replay cannot assign the"110 "requested recorded address (%p, %p)",111 VAddr, MemoryStart);112 }113 114 INFO(OMP_INFOTYPE_PLUGIN_KERNEL, Device->getDeviceId(),115 "Allocated %" PRIu64 " bytes at %p for replay.\n", ASize, MemoryStart);116 117 MemoryPtr = MemoryStart;118 MemorySize = 0;119 TotalSize = ASize;120 UsedVAMap = true;121 return Plugin::success();122 }123 124 Error preAllocateHeuristic(uint64_t MaxMemoryAllocation,125 uint64_t RequiredMemoryAllocation, void *VAddr) {126 const size_t MAX_MEMORY_ALLOCATION = MaxMemoryAllocation;127 constexpr size_t STEP = 1024 * 1024 * 1024ULL;128 MemoryStart = nullptr;129 for (TotalSize = MAX_MEMORY_ALLOCATION; TotalSize > 0; TotalSize -= STEP) {130 auto MemoryStartOrErr =131 Device->allocate(TotalSize, /*HstPtr=*/nullptr, TARGET_ALLOC_DEFAULT);132 if (!MemoryStartOrErr)133 return MemoryStartOrErr.takeError();134 MemoryStart = *MemoryStartOrErr;135 if (MemoryStart)136 break;137 }138 if (!MemoryStart)139 return Plugin::error(ErrorCode::INVALID_ARGUMENT,140 "allocating record/replay memory");141 142 if (VAddr && VAddr != MemoryStart)143 MemoryOffset = uintptr_t(VAddr) - uintptr_t(MemoryStart);144 145 MemoryPtr = MemoryStart;146 MemorySize = 0;147 148 // Check if we need adjustment.149 if (MemoryOffset > 0 &&150 TotalSize >= RequiredMemoryAllocation + MemoryOffset) {151 // If we are off but "before" the required address and with enough space,152 // we just "allocate" the offset to match the required address.153 MemoryPtr = (char *)MemoryPtr + MemoryOffset;154 MemorySize += MemoryOffset;155 MemoryOffset = 0;156 assert(MemoryPtr == VAddr && "Expected offset adjustment to work");157 } else if (MemoryOffset) {158 // If we are off and in a situation we cannot just "waste" memory to force159 // a match, we hope adjusting the arguments is sufficient.160 REPORT(161 "WARNING Failed to allocate replay memory at required location %p, "162 "got %p, trying to offset argument pointers by %" PRIi64 "\n",163 VAddr, MemoryStart, MemoryOffset);164 }165 166 INFO(OMP_INFOTYPE_PLUGIN_KERNEL, Device->getDeviceId(),167 "Allocated %" PRIu64 " bytes at %p for replay.\n", TotalSize,168 MemoryStart);169 170 return Plugin::success();171 }172 173 Error preallocateDeviceMemory(uint64_t DeviceMemorySize, void *ReqVAddr) {174 if (Device->supportVAManagement()) {175 auto Err = preAllocateVAMemory(DeviceMemorySize, ReqVAddr);176 if (Err) {177 REPORT("WARNING VA mapping failed, fallback to heuristic: "178 "(Error: %s)\n",179 toString(std::move(Err)).data());180 }181 }182 183 uint64_t DevMemSize;184 if (Device->getDeviceMemorySize(DevMemSize))185 return Plugin::error(ErrorCode::UNKNOWN,186 "cannot determine Device Memory Size");187 188 return preAllocateHeuristic(DevMemSize, DeviceMemorySize, ReqVAddr);189 }190 191 void dumpDeviceMemory(StringRef Filename) {192 ErrorOr<std::unique_ptr<WritableMemoryBuffer>> DeviceMemoryMB =193 WritableMemoryBuffer::getNewUninitMemBuffer(MemorySize);194 if (!DeviceMemoryMB)195 report_fatal_error("Error creating MemoryBuffer for device memory");196 197 auto Err = Device->dataRetrieve(DeviceMemoryMB.get()->getBufferStart(),198 MemoryStart, MemorySize, nullptr);199 if (Err)200 report_fatal_error("Error retrieving data for target pointer");201 202 StringRef DeviceMemory(DeviceMemoryMB.get()->getBufferStart(), MemorySize);203 std::error_code EC;204 raw_fd_ostream OS(Filename, EC);205 if (EC)206 report_fatal_error("Error dumping memory to file " + Filename + " :" +207 EC.message());208 OS << DeviceMemory;209 OS.close();210 }211 212public:213 bool isRecording() const { return Status == RRStatusTy::RRRecording; }214 bool isReplaying() const { return Status == RRStatusTy::RRReplaying; }215 bool isRecordingOrReplaying() const {216 return (Status != RRStatusTy::RRDeactivated);217 }218 void setStatus(RRStatusTy Status) { this->Status = Status; }219 bool isSaveOutputEnabled() const { return ReplaySaveOutput; }220 void addEntry(const char *Name, uint64_t Size, void *Addr) {221 GlobalEntries.emplace_back(GlobalEntry{Name, Size, Addr});222 }223 224 void saveImage(const char *Name, const DeviceImageTy &Image) {225 SmallString<128> ImageName = {Name, ".image"};226 std::error_code EC;227 raw_fd_ostream OS(ImageName, EC);228 if (EC)229 report_fatal_error("Error saving image : " + StringRef(EC.message()));230 OS << Image.getMemoryBuffer().getBuffer();231 OS.close();232 }233 234 void dumpGlobals(StringRef Filename, DeviceImageTy &Image) {235 int32_t Size = 0;236 237 for (auto &OffloadEntry : GlobalEntries) {238 if (!OffloadEntry.Size)239 continue;240 // Get the total size of the string and entry including the null byte.241 Size += std::strlen(OffloadEntry.Name) + 1 + sizeof(uint32_t) +242 OffloadEntry.Size;243 }244 245 ErrorOr<std::unique_ptr<WritableMemoryBuffer>> GlobalsMB =246 WritableMemoryBuffer::getNewUninitMemBuffer(Size);247 if (!GlobalsMB)248 report_fatal_error("Error creating MemoryBuffer for globals memory");249 250 void *BufferPtr = GlobalsMB.get()->getBufferStart();251 for (auto &OffloadEntry : GlobalEntries) {252 if (!OffloadEntry.Size)253 continue;254 255 int32_t NameLength = std::strlen(OffloadEntry.Name) + 1;256 memcpy(BufferPtr, OffloadEntry.Name, NameLength);257 BufferPtr = utils::advancePtr(BufferPtr, NameLength);258 259 *((uint32_t *)(BufferPtr)) = OffloadEntry.Size;260 BufferPtr = utils::advancePtr(BufferPtr, sizeof(uint32_t));261 262 auto Err = Plugin::success();263 {264 if (auto Err = Device->dataRetrieve(BufferPtr, OffloadEntry.Addr,265 OffloadEntry.Size, nullptr))266 report_fatal_error("Error retrieving data for global");267 }268 if (Err)269 report_fatal_error("Error retrieving data for global");270 BufferPtr = utils::advancePtr(BufferPtr, OffloadEntry.Size);271 }272 assert(BufferPtr == GlobalsMB->get()->getBufferEnd() &&273 "Buffer over/under-filled.");274 assert(Size == utils::getPtrDiff(BufferPtr,275 GlobalsMB->get()->getBufferStart()) &&276 "Buffer size mismatch");277 278 StringRef GlobalsMemory(GlobalsMB.get()->getBufferStart(), Size);279 std::error_code EC;280 raw_fd_ostream OS(Filename, EC);281 OS << GlobalsMemory;282 OS.close();283 }284 285 void saveKernelDescr(const char *Name, KernelLaunchParamsTy LaunchParams,286 int32_t NumArgs, uint64_t NumTeamsClause,287 uint32_t ThreadLimitClause, uint64_t LoopTripCount) {288 json::Object JsonKernelInfo;289 JsonKernelInfo["Name"] = Name;290 JsonKernelInfo["NumArgs"] = NumArgs;291 JsonKernelInfo["NumTeamsClause"] = NumTeamsClause;292 JsonKernelInfo["ThreadLimitClause"] = ThreadLimitClause;293 JsonKernelInfo["LoopTripCount"] = LoopTripCount;294 JsonKernelInfo["DeviceMemorySize"] = MemorySize;295 JsonKernelInfo["DeviceId"] = Device->getDeviceId();296 JsonKernelInfo["BumpAllocVAStart"] = (intptr_t)MemoryStart;297 298 json::Array JsonArgPtrs;299 for (int I = 0; I < NumArgs; ++I)300 JsonArgPtrs.push_back((intptr_t)LaunchParams.Ptrs[I]);301 JsonKernelInfo["ArgPtrs"] = json::Value(std::move(JsonArgPtrs));302 303 json::Array JsonArgOffsets;304 for (int I = 0; I < NumArgs; ++I)305 JsonArgOffsets.push_back(0);306 JsonKernelInfo["ArgOffsets"] = json::Value(std::move(JsonArgOffsets));307 308 SmallString<128> JsonFilename = {Name, ".json"};309 std::error_code EC;310 raw_fd_ostream JsonOS(JsonFilename.str(), EC);311 if (EC)312 report_fatal_error("Error saving kernel json file : " +313 StringRef(EC.message()));314 JsonOS << json::Value(std::move(JsonKernelInfo));315 JsonOS.close();316 }317 318 void saveKernelInput(const char *Name, DeviceImageTy &Image) {319 SmallString<128> GlobalsFilename = {Name, ".globals"};320 dumpGlobals(GlobalsFilename, Image);321 322 SmallString<128> MemoryFilename = {Name, ".memory"};323 dumpDeviceMemory(MemoryFilename);324 }325 326 void saveKernelOutputInfo(const char *Name) {327 SmallString<128> OutputFilename = {328 Name, (isRecording() ? ".original.output" : ".replay.output")};329 dumpDeviceMemory(OutputFilename);330 }331 332 void *alloc(uint64_t Size) {333 assert(MemoryStart && "Expected memory has been pre-allocated");334 void *Alloc = nullptr;335 constexpr int Alignment = 16;336 // Assumes alignment is a power of 2.337 int64_t AlignedSize = (Size + (Alignment - 1)) & (~(Alignment - 1));338 std::lock_guard<std::mutex> LG(AllocationLock);339 Alloc = MemoryPtr;340 MemoryPtr = (char *)MemoryPtr + AlignedSize;341 MemorySize += AlignedSize;342 DP("Memory Allocator return " DPxMOD "\n", DPxPTR(Alloc));343 return Alloc;344 }345 346 Error init(GenericDeviceTy *Device, uint64_t MemSize, void *VAddr,347 RRStatusTy Status, bool SaveOutput, uint64_t &ReqPtrArgOffset) {348 this->Device = Device;349 this->Status = Status;350 this->ReplaySaveOutput = SaveOutput;351 352 if (auto Err = preallocateDeviceMemory(MemSize, VAddr))353 return Err;354 355 INFO(OMP_INFOTYPE_PLUGIN_KERNEL, Device->getDeviceId(),356 "Record Replay Initialized (%p)"357 " as starting address, %lu Memory Size"358 " and set on status %s\n",359 MemoryStart, TotalSize,360 Status == RRStatusTy::RRRecording ? "Recording" : "Replaying");361 362 // Tell the user to offset pointer arguments as the memory allocation does363 // not match.364 ReqPtrArgOffset = MemoryOffset;365 return Plugin::success();366 }367 368 Error deinit() {369 if (UsedVAMap) {370 if (auto Err = Device->memoryVAUnMap(MemoryStart, TotalSize))371 return Err;372 } else {373 if (auto Err = Device->free(MemoryStart))374 return Err;375 }376 return Plugin::success();377 }378};379} // namespace llvm::omp::target::plugin380 381AsyncInfoWrapperTy::AsyncInfoWrapperTy(GenericDeviceTy &Device,382 __tgt_async_info *AsyncInfoPtr)383 : Device(Device),384 AsyncInfoPtr(AsyncInfoPtr ? AsyncInfoPtr : &LocalAsyncInfo) {}385 386void AsyncInfoWrapperTy::finalize(Error &Err) {387 assert(AsyncInfoPtr && "AsyncInfoWrapperTy already finalized");388 389 // If we used a local async info object we want synchronous behavior. In that390 // case, and assuming the current status code is correct, we will synchronize391 // explicitly when the object is deleted. Update the error with the result of392 // the synchronize operation.393 if (AsyncInfoPtr == &LocalAsyncInfo && LocalAsyncInfo.Queue && !Err)394 Err = Device.synchronize(&LocalAsyncInfo);395 396 // Invalidate the wrapper object.397 AsyncInfoPtr = nullptr;398}399 400Error GenericKernelTy::init(GenericDeviceTy &GenericDevice,401 DeviceImageTy &Image) {402 403 ImagePtr = &Image;404 405 // Retrieve kernel environment object for the kernel.406 std::string EnvironmentName = std::string(Name) + "_kernel_environment";407 GenericGlobalHandlerTy &GHandler = GenericDevice.Plugin.getGlobalHandler();408 if (GHandler.isSymbolInImage(GenericDevice, Image, EnvironmentName)) {409 GlobalTy KernelEnv(EnvironmentName, sizeof(KernelEnvironment),410 &KernelEnvironment);411 if (auto Err =412 GHandler.readGlobalFromImage(GenericDevice, *ImagePtr, KernelEnv))413 return Err;414 } else {415 KernelEnvironment = KernelEnvironmentTy{};416 DP("Failed to read kernel environment for '%s' Using default Bare (0) "417 "execution mode\n",418 getName());419 }420 421 // Max = Config.Max > 0 ? min(Config.Max, Device.Max) : Device.Max;422 MaxNumThreads = KernelEnvironment.Configuration.MaxThreads > 0423 ? std::min(KernelEnvironment.Configuration.MaxThreads,424 int32_t(GenericDevice.getThreadLimit()))425 : GenericDevice.getThreadLimit();426 427 // Pref = Config.Pref > 0 ? max(Config.Pref, Device.Pref) : Device.Pref;428 PreferredNumThreads =429 KernelEnvironment.Configuration.MinThreads > 0430 ? std::max(KernelEnvironment.Configuration.MinThreads,431 int32_t(GenericDevice.getDefaultNumThreads()))432 : GenericDevice.getDefaultNumThreads();433 434 return initImpl(GenericDevice, Image);435}436 437Expected<KernelLaunchEnvironmentTy *>438GenericKernelTy::getKernelLaunchEnvironment(439 GenericDeviceTy &GenericDevice, uint32_t Version,440 AsyncInfoWrapperTy &AsyncInfoWrapper) const {441 // Ctor/Dtor have no arguments, replaying uses the original kernel launch442 // environment. Older versions of the compiler do not generate a kernel443 // launch environment.444 if (GenericDevice.Plugin.getRecordReplay().isReplaying() ||445 Version < OMP_KERNEL_ARG_MIN_VERSION_WITH_DYN_PTR)446 return nullptr;447 448 if (!KernelEnvironment.Configuration.ReductionDataSize ||449 !KernelEnvironment.Configuration.ReductionBufferLength)450 return reinterpret_cast<KernelLaunchEnvironmentTy *>(~0);451 452 // TODO: Check if the kernel needs a launch environment.453 auto AllocOrErr = GenericDevice.dataAlloc(sizeof(KernelLaunchEnvironmentTy),454 /*HostPtr=*/nullptr,455 TargetAllocTy::TARGET_ALLOC_DEVICE);456 if (!AllocOrErr)457 return AllocOrErr.takeError();458 459 // Remember to free the memory later.460 AsyncInfoWrapper.freeAllocationAfterSynchronization(*AllocOrErr);461 462 /// Use the KLE in the __tgt_async_info to ensure a stable address for the463 /// async data transfer.464 auto &LocalKLE = (*AsyncInfoWrapper).KernelLaunchEnvironment;465 LocalKLE = KernelLaunchEnvironment;466 {467 auto AllocOrErr = GenericDevice.dataAlloc(468 KernelEnvironment.Configuration.ReductionDataSize *469 KernelEnvironment.Configuration.ReductionBufferLength,470 /*HostPtr=*/nullptr, TargetAllocTy::TARGET_ALLOC_DEVICE);471 if (!AllocOrErr)472 return AllocOrErr.takeError();473 LocalKLE.ReductionBuffer = *AllocOrErr;474 // Remember to free the memory later.475 AsyncInfoWrapper.freeAllocationAfterSynchronization(*AllocOrErr);476 }477 478 INFO(OMP_INFOTYPE_DATA_TRANSFER, GenericDevice.getDeviceId(),479 "Copying data from host to device, HstPtr=" DPxMOD ", TgtPtr=" DPxMOD480 ", Size=%" PRId64 ", Name=KernelLaunchEnv\n",481 DPxPTR(&LocalKLE), DPxPTR(*AllocOrErr),482 sizeof(KernelLaunchEnvironmentTy));483 484 auto Err = GenericDevice.dataSubmit(*AllocOrErr, &LocalKLE,485 sizeof(KernelLaunchEnvironmentTy),486 AsyncInfoWrapper);487 if (Err)488 return Err;489 return static_cast<KernelLaunchEnvironmentTy *>(*AllocOrErr);490}491 492Error GenericKernelTy::printLaunchInfo(GenericDeviceTy &GenericDevice,493 KernelArgsTy &KernelArgs,494 uint32_t NumThreads[3],495 uint32_t NumBlocks[3]) const {496 INFO(OMP_INFOTYPE_PLUGIN_KERNEL, GenericDevice.getDeviceId(),497 "Launching kernel %s with [%u,%u,%u] blocks and [%u,%u,%u] threads in "498 "%s mode\n",499 getName(), NumBlocks[0], NumBlocks[1], NumBlocks[2], NumThreads[0],500 NumThreads[1], NumThreads[2], getExecutionModeName());501 return printLaunchInfoDetails(GenericDevice, KernelArgs, NumThreads,502 NumBlocks);503}504 505Error GenericKernelTy::printLaunchInfoDetails(GenericDeviceTy &GenericDevice,506 KernelArgsTy &KernelArgs,507 uint32_t NumThreads[3],508 uint32_t NumBlocks[3]) const {509 return Plugin::success();510}511 512Error GenericKernelTy::launch(GenericDeviceTy &GenericDevice, void **ArgPtrs,513 ptrdiff_t *ArgOffsets, KernelArgsTy &KernelArgs,514 AsyncInfoWrapperTy &AsyncInfoWrapper) const {515 llvm::SmallVector<void *, 16> Args;516 llvm::SmallVector<void *, 16> Ptrs;517 518 auto KernelLaunchEnvOrErr = getKernelLaunchEnvironment(519 GenericDevice, KernelArgs.Version, AsyncInfoWrapper);520 if (!KernelLaunchEnvOrErr)521 return KernelLaunchEnvOrErr.takeError();522 523 KernelLaunchParamsTy LaunchParams;524 525 // Kernel languages don't use indirection.526 if (KernelArgs.Flags.IsCUDA) {527 LaunchParams =528 *reinterpret_cast<KernelLaunchParamsTy *>(KernelArgs.ArgPtrs);529 } else {530 LaunchParams =531 prepareArgs(GenericDevice, ArgPtrs, ArgOffsets, KernelArgs.NumArgs,532 Args, Ptrs, *KernelLaunchEnvOrErr);533 }534 535 uint32_t NumThreads[3] = {KernelArgs.ThreadLimit[0],536 KernelArgs.ThreadLimit[1],537 KernelArgs.ThreadLimit[2]};538 uint32_t NumBlocks[3] = {KernelArgs.NumTeams[0], KernelArgs.NumTeams[1],539 KernelArgs.NumTeams[2]};540 if (!isBareMode()) {541 NumThreads[0] = getNumThreads(GenericDevice, NumThreads);542 NumBlocks[0] = getNumBlocks(GenericDevice, NumBlocks, KernelArgs.Tripcount,543 NumThreads[0], KernelArgs.ThreadLimit[0] > 0);544 }545 546 // Record the kernel description after we modified the argument count and num547 // blocks/threads.548 RecordReplayTy &RecordReplay = GenericDevice.Plugin.getRecordReplay();549 if (RecordReplay.isRecording()) {550 RecordReplay.saveImage(getName(), getImage());551 RecordReplay.saveKernelInput(getName(), getImage());552 RecordReplay.saveKernelDescr(getName(), LaunchParams, KernelArgs.NumArgs,553 NumBlocks[0], NumThreads[0],554 KernelArgs.Tripcount);555 }556 557 if (auto Err =558 printLaunchInfo(GenericDevice, KernelArgs, NumThreads, NumBlocks))559 return Err;560 561 return launchImpl(GenericDevice, NumThreads, NumBlocks, KernelArgs,562 LaunchParams, AsyncInfoWrapper);563}564 565KernelLaunchParamsTy GenericKernelTy::prepareArgs(566 GenericDeviceTy &GenericDevice, void **ArgPtrs, ptrdiff_t *ArgOffsets,567 uint32_t &NumArgs, llvm::SmallVectorImpl<void *> &Args,568 llvm::SmallVectorImpl<void *> &Ptrs,569 KernelLaunchEnvironmentTy *KernelLaunchEnvironment) const {570 uint32_t KLEOffset = !!KernelLaunchEnvironment;571 NumArgs += KLEOffset;572 573 if (NumArgs == 0)574 return KernelLaunchParamsTy{};575 576 Args.resize(NumArgs);577 Ptrs.resize(NumArgs);578 579 if (KernelLaunchEnvironment) {580 Args[0] = KernelLaunchEnvironment;581 Ptrs[0] = &Args[0];582 }583 584 for (uint32_t I = KLEOffset; I < NumArgs; ++I) {585 Args[I] =586 (void *)((intptr_t)ArgPtrs[I - KLEOffset] + ArgOffsets[I - KLEOffset]);587 Ptrs[I] = &Args[I];588 }589 return KernelLaunchParamsTy{sizeof(void *) * NumArgs, &Args[0], &Ptrs[0]};590}591 592uint32_t GenericKernelTy::getNumThreads(GenericDeviceTy &GenericDevice,593 uint32_t ThreadLimitClause[3]) const {594 assert(!isBareMode() && "bare kernel should not call this function");595 596 assert(ThreadLimitClause[1] == 1 && ThreadLimitClause[2] == 1 &&597 "Multi dimensional launch not supported yet.");598 599 if (ThreadLimitClause[0] > 0 && isGenericMode())600 ThreadLimitClause[0] += GenericDevice.getWarpSize();601 602 return std::min(MaxNumThreads, (ThreadLimitClause[0] > 0)603 ? ThreadLimitClause[0]604 : PreferredNumThreads);605}606 607uint32_t GenericKernelTy::getNumBlocks(GenericDeviceTy &GenericDevice,608 uint32_t NumTeamsClause[3],609 uint64_t LoopTripCount,610 uint32_t &NumThreads,611 bool IsNumThreadsFromUser) const {612 assert(!isBareMode() && "bare kernel should not call this function");613 614 assert(NumTeamsClause[1] == 1 && NumTeamsClause[2] == 1 &&615 "Multi dimensional launch not supported yet.");616 617 if (NumTeamsClause[0] > 0) {618 // TODO: We need to honor any value and consequently allow more than the619 // block limit. For this we might need to start multiple kernels or let the620 // blocks start again until the requested number has been started.621 return std::min(NumTeamsClause[0], GenericDevice.getBlockLimit());622 }623 624 // Return the number of teams required to cover the loop iterations.625 if (isNoLoopMode())626 return LoopTripCount > 0 ? (((LoopTripCount - 1) / NumThreads) + 1) : 1;627 628 uint64_t DefaultNumBlocks = GenericDevice.getDefaultNumBlocks();629 uint64_t TripCountNumBlocks = std::numeric_limits<uint64_t>::max();630 if (LoopTripCount > 0) {631 if (isSPMDMode()) {632 // We have a combined construct, i.e. `target teams distribute633 // parallel for [simd]`. We launch so many teams so that each thread634 // will execute one iteration of the loop; rounded up to the nearest635 // integer. However, if that results in too few teams, we artificially636 // reduce the thread count per team to increase the outer parallelism.637 auto MinThreads = GenericDevice.getMinThreadsForLowTripCountLoop();638 MinThreads = std::min(MinThreads, NumThreads);639 640 // Honor the thread_limit clause; only lower the number of threads.641 [[maybe_unused]] auto OldNumThreads = NumThreads;642 if (LoopTripCount >= DefaultNumBlocks * NumThreads ||643 IsNumThreadsFromUser) {644 // Enough parallelism for teams and threads.645 TripCountNumBlocks = ((LoopTripCount - 1) / NumThreads) + 1;646 assert(IsNumThreadsFromUser ||647 TripCountNumBlocks >= DefaultNumBlocks &&648 "Expected sufficient outer parallelism.");649 } else if (LoopTripCount >= DefaultNumBlocks * MinThreads) {650 // Enough parallelism for teams, limit threads.651 652 // This case is hard; for now, we force "full warps":653 // First, compute a thread count assuming DefaultNumBlocks.654 auto NumThreadsDefaultBlocks =655 (LoopTripCount + DefaultNumBlocks - 1) / DefaultNumBlocks;656 // Now get a power of two that is larger or equal.657 auto NumThreadsDefaultBlocksP2 =658 llvm::PowerOf2Ceil(NumThreadsDefaultBlocks);659 // Do not increase a thread limit given be the user.660 NumThreads = std::min(NumThreads, uint32_t(NumThreadsDefaultBlocksP2));661 assert(NumThreads >= MinThreads &&662 "Expected sufficient inner parallelism.");663 TripCountNumBlocks = ((LoopTripCount - 1) / NumThreads) + 1;664 } else {665 // Not enough parallelism for teams and threads, limit both.666 NumThreads = std::min(NumThreads, MinThreads);667 TripCountNumBlocks = ((LoopTripCount - 1) / NumThreads) + 1;668 }669 670 assert(NumThreads * TripCountNumBlocks >= LoopTripCount &&671 "Expected sufficient parallelism");672 assert(OldNumThreads >= NumThreads &&673 "Number of threads cannot be increased!");674 } else {675 assert((isGenericMode() || isGenericSPMDMode()) &&676 "Unexpected execution mode!");677 // If we reach this point, then we have a non-combined construct, i.e.678 // `teams distribute` with a nested `parallel for` and each team is679 // assigned one iteration of the `distribute` loop. E.g.:680 //681 // #pragma omp target teams distribute682 // for(...loop_tripcount...) {683 // #pragma omp parallel for684 // for(...) {}685 // }686 //687 // Threads within a team will execute the iterations of the `parallel`688 // loop.689 TripCountNumBlocks = LoopTripCount;690 }691 }692 693 uint32_t PreferredNumBlocks = TripCountNumBlocks;694 // If the loops are long running we rather reuse blocks than spawn too many.695 if (GenericDevice.getReuseBlocksForHighTripCount())696 PreferredNumBlocks = std::min(TripCountNumBlocks, DefaultNumBlocks);697 return std::min(PreferredNumBlocks, GenericDevice.getBlockLimit());698}699 700GenericDeviceTy::GenericDeviceTy(GenericPluginTy &Plugin, int32_t DeviceId,701 int32_t NumDevices,702 const llvm::omp::GV &OMPGridValues)703 : Plugin(Plugin), MemoryManager(nullptr), OMP_TeamLimit("OMP_TEAM_LIMIT"),704 OMP_NumTeams("OMP_NUM_TEAMS"),705 OMP_TeamsThreadLimit("OMP_TEAMS_THREAD_LIMIT"),706 OMPX_DebugKind("LIBOMPTARGET_DEVICE_RTL_DEBUG"),707 OMPX_SharedMemorySize("LIBOMPTARGET_SHARED_MEMORY_SIZE"),708 // Do not initialize the following two envars since they depend on the709 // device initialization. These cannot be consulted until the device is710 // initialized correctly. We initialize them in GenericDeviceTy::init().711 OMPX_TargetStackSize(), OMPX_TargetHeapSize(),712 // By default, the initial number of streams and events is 1.713 OMPX_InitialNumStreams("LIBOMPTARGET_NUM_INITIAL_STREAMS", 1),714 OMPX_InitialNumEvents("LIBOMPTARGET_NUM_INITIAL_EVENTS", 1),715 DeviceId(DeviceId), GridValues(OMPGridValues),716 PeerAccesses(NumDevices, PeerAccessState::PENDING), PeerAccessesLock(),717 PinnedAllocs(*this), RPCServer(nullptr) {718 // Conservative fall-back to the plugin's device uid for the case that no real719 // vendor (u)uid will become available later.720 setDeviceUidFromVendorUid(std::to_string(static_cast<uint64_t>(DeviceId)));721 722#ifdef OMPT_SUPPORT723 OmptInitialized.store(false);724 // Bind the callbacks to this device's member functions725#define bindOmptCallback(Name, Type, Code) \726 if (ompt::Initialized && ompt::lookupCallbackByCode) { \727 ompt::lookupCallbackByCode((ompt_callbacks_t)(Code), \728 ((ompt_callback_t *)&(Name##_fn))); \729 DP("OMPT: class bound %s=%p\n", #Name, ((void *)(uint64_t)Name##_fn)); \730 }731 732 FOREACH_OMPT_DEVICE_EVENT(bindOmptCallback);733#undef bindOmptCallback734 735#endif736}737 738Error GenericDeviceTy::init(GenericPluginTy &Plugin) {739 if (auto Err = initImpl(Plugin))740 return Err;741 742#ifdef OMPT_SUPPORT743 if (ompt::Initialized) {744 bool ExpectedStatus = false;745 if (OmptInitialized.compare_exchange_strong(ExpectedStatus, true))746 performOmptCallback(device_initialize, Plugin.getUserId(DeviceId),747 /*type=*/getComputeUnitKind().c_str(),748 /*device=*/reinterpret_cast<ompt_device_t *>(this),749 /*lookup=*/ompt::lookupCallbackByName,750 /*documentation=*/nullptr);751 }752#endif753 754 // Read and reinitialize the envars that depend on the device initialization.755 // Notice these two envars may change the stack size and heap size of the756 // device, so they need the device properly initialized.757 auto StackSizeEnvarOrErr = UInt64Envar::create(758 "LIBOMPTARGET_STACK_SIZE",759 [this](uint64_t &V) -> Error { return getDeviceStackSize(V); },760 [this](uint64_t V) -> Error { return setDeviceStackSize(V); });761 if (!StackSizeEnvarOrErr)762 return StackSizeEnvarOrErr.takeError();763 OMPX_TargetStackSize = std::move(*StackSizeEnvarOrErr);764 765 if (hasDeviceHeapSize()) {766 auto HeapSizeEnvarOrErr = UInt64Envar::create(767 "LIBOMPTARGET_HEAP_SIZE",768 [this](uint64_t &V) -> Error { return getDeviceHeapSize(V); },769 [this](uint64_t V) -> Error { return setDeviceHeapSize(V); });770 if (!HeapSizeEnvarOrErr)771 return HeapSizeEnvarOrErr.takeError();772 OMPX_TargetHeapSize = std::move(*HeapSizeEnvarOrErr);773 }774 775 // Update the maximum number of teams and threads after the device776 // initialization sets the corresponding hardware limit.777 if (OMP_NumTeams > 0)778 GridValues.GV_Max_Teams =779 std::min(GridValues.GV_Max_Teams, uint32_t(OMP_NumTeams));780 781 if (OMP_TeamsThreadLimit > 0)782 GridValues.GV_Max_WG_Size =783 std::min(GridValues.GV_Max_WG_Size, uint32_t(OMP_TeamsThreadLimit));784 785 // Enable the memory manager if required.786 auto [ThresholdMM, EnableMM] = MemoryManagerTy::getSizeThresholdFromEnv();787 if (EnableMM) {788 if (ThresholdMM == 0)789 ThresholdMM = getMemoryManagerSizeThreshold();790 MemoryManager = new MemoryManagerTy(*this, ThresholdMM);791 }792 793 return Plugin::success();794}795 796Error GenericDeviceTy::unloadBinary(DeviceImageTy *Image) {797 if (auto Err = callGlobalDestructors(Plugin, *Image))798 return Err;799 800 GenericGlobalHandlerTy &Handler = Plugin.getGlobalHandler();801 auto ProfOrErr = Handler.readProfilingGlobals(*this, *Image);802 if (!ProfOrErr)803 return ProfOrErr.takeError();804 805 if (!ProfOrErr->empty()) {806 // Dump out profdata807 if ((OMPX_DebugKind.get() & uint32_t(DeviceDebugKind::PGODump)) ==808 uint32_t(DeviceDebugKind::PGODump))809 ProfOrErr->dump();810 811 // Write data to profiling file812 if (auto Err = ProfOrErr->write())813 return Err;814 }815 816 return unloadBinaryImpl(Image);817}818 819Error GenericDeviceTy::deinit(GenericPluginTy &Plugin) {820 for (auto &I : LoadedImages)821 if (auto Err = unloadBinary(I))822 return Err;823 LoadedImages.clear();824 825 // Delete the memory manager before deinitializing the device. Otherwise,826 // we may delete device allocations after the device is deinitialized.827 if (MemoryManager)828 delete MemoryManager;829 MemoryManager = nullptr;830 831 RecordReplayTy &RecordReplay = Plugin.getRecordReplay();832 if (RecordReplay.isRecordingOrReplaying())833 if (auto Err = RecordReplay.deinit())834 return Err;835 836 if (RPCServer)837 if (auto Err = RPCServer->deinitDevice(*this))838 return Err;839 840#ifdef OMPT_SUPPORT841 if (ompt::Initialized) {842 bool ExpectedStatus = true;843 if (OmptInitialized.compare_exchange_strong(ExpectedStatus, false))844 performOmptCallback(device_finalize, Plugin.getUserId(DeviceId));845 }846#endif847 848 return deinitImpl();849}850Expected<DeviceImageTy *> GenericDeviceTy::loadBinary(GenericPluginTy &Plugin,851 StringRef InputTgtImage) {852 DP("Load data from image " DPxMOD "\n", DPxPTR(InputTgtImage.bytes_begin()));853 854 std::unique_ptr<MemoryBuffer> Buffer;855 if (identify_magic(InputTgtImage) == file_magic::bitcode) {856 auto CompiledImageOrErr = Plugin.getJIT().process(InputTgtImage, *this);857 if (!CompiledImageOrErr) {858 return Plugin::error(ErrorCode::COMPILE_FAILURE,859 CompiledImageOrErr.takeError(),860 "failure to jit IR image");861 }862 Buffer = std::move(*CompiledImageOrErr);863 } else {864 Buffer = MemoryBuffer::getMemBufferCopy(InputTgtImage);865 }866 867 // Load the binary and allocate the image object. Use the next available id868 // for the image id, which is the number of previously loaded images.869 auto ImageOrErr = loadBinaryImpl(std::move(Buffer), LoadedImages.size());870 if (!ImageOrErr)871 return ImageOrErr.takeError();872 DeviceImageTy *Image = *ImageOrErr;873 874 // Add the image to list.875 LoadedImages.push_back(Image);876 877 if (auto Err = setupRPCServer(Plugin, *Image))878 return std::move(Err);879 880#ifdef OMPT_SUPPORT881 if (ompt::Initialized) {882 size_t Bytes = InputTgtImage.size();883 performOmptCallback(884 device_load, Plugin.getUserId(DeviceId),885 /*FileName=*/nullptr, /*FileOffset=*/0, /*VmaInFile=*/nullptr,886 /*ImgSize=*/Bytes,887 /*HostAddr=*/const_cast<unsigned char *>(InputTgtImage.bytes_begin()),888 /*DeviceAddr=*/nullptr, /* FIXME: ModuleId */ 0);889 }890#endif891 892 // Call any global constructors present on the device.893 if (auto Err = callGlobalConstructors(Plugin, *Image))894 return std::move(Err);895 896 // Return the pointer to the table of entries.897 return Image;898}899 900Error GenericDeviceTy::setupRPCServer(GenericPluginTy &Plugin,901 DeviceImageTy &Image) {902 // The plugin either does not need an RPC server or it is unavailable.903 if (!shouldSetupRPCServer())904 return Plugin::success();905 906 // Check if this device needs to run an RPC server.907 RPCServerTy &Server = Plugin.getRPCServer();908 auto UsingOrErr =909 Server.isDeviceUsingRPC(*this, Plugin.getGlobalHandler(), Image);910 if (!UsingOrErr)911 return UsingOrErr.takeError();912 913 if (!UsingOrErr.get())914 return Plugin::success();915 916 if (auto Err = Server.initDevice(*this, Plugin.getGlobalHandler(), Image))917 return Err;918 919 if (auto Err = Server.startThread())920 return Err;921 922 RPCServer = &Server;923 DP("Running an RPC server on device %d\n", getDeviceId());924 return Plugin::success();925}926 927Error PinnedAllocationMapTy::insertEntry(void *HstPtr, void *DevAccessiblePtr,928 size_t Size, bool ExternallyLocked) {929 // Insert the new entry into the map.930 auto Res = Allocs.insert({HstPtr, DevAccessiblePtr, Size, ExternallyLocked});931 if (!Res.second)932 return Plugin::error(ErrorCode::INVALID_ARGUMENT,933 "cannot insert locked buffer entry");934 935 // Check whether the next entry overlaps with the inserted entry.936 auto It = std::next(Res.first);937 if (It == Allocs.end())938 return Plugin::success();939 940 const EntryTy *NextEntry = &(*It);941 if (intersects(NextEntry->HstPtr, NextEntry->Size, HstPtr, Size))942 return Plugin::error(ErrorCode::INVALID_ARGUMENT,943 "partial overlapping not allowed in locked buffers");944 945 return Plugin::success();946}947 948Error PinnedAllocationMapTy::eraseEntry(const EntryTy &Entry) {949 // Erase the existing entry. Notice this requires an additional map lookup,950 // but this should not be a performance issue. Using iterators would make951 // the code more difficult to read.952 size_t Erased = Allocs.erase({Entry.HstPtr});953 if (!Erased)954 return Plugin::error(ErrorCode::INVALID_ARGUMENT,955 "cannot erase locked buffer entry");956 return Plugin::success();957}958 959Error PinnedAllocationMapTy::registerEntryUse(const EntryTy &Entry,960 void *HstPtr, size_t Size) {961 if (!contains(Entry.HstPtr, Entry.Size, HstPtr, Size))962 return Plugin::error(ErrorCode::INVALID_ARGUMENT,963 "partial overlapping not allowed in locked buffers");964 965 ++Entry.References;966 return Plugin::success();967}968 969Expected<bool> PinnedAllocationMapTy::unregisterEntryUse(const EntryTy &Entry) {970 if (Entry.References == 0)971 return Plugin::error(ErrorCode::INVALID_ARGUMENT,972 "invalid number of references");973 974 // Return whether this was the last user.975 return (--Entry.References == 0);976}977 978Error PinnedAllocationMapTy::registerHostBuffer(void *HstPtr,979 void *DevAccessiblePtr,980 size_t Size) {981 assert(HstPtr && "Invalid pointer");982 assert(DevAccessiblePtr && "Invalid pointer");983 assert(Size && "Invalid size");984 985 std::lock_guard<std::shared_mutex> Lock(Mutex);986 987 // No pinned allocation should intersect.988 const EntryTy *Entry = findIntersecting(HstPtr);989 if (Entry)990 return Plugin::error(ErrorCode::INVALID_ARGUMENT,991 "cannot insert entry due to an existing one");992 993 // Now insert the new entry.994 return insertEntry(HstPtr, DevAccessiblePtr, Size);995}996 997Error PinnedAllocationMapTy::unregisterHostBuffer(void *HstPtr) {998 assert(HstPtr && "Invalid pointer");999 1000 std::lock_guard<std::shared_mutex> Lock(Mutex);1001 1002 const EntryTy *Entry = findIntersecting(HstPtr);1003 if (!Entry)1004 return Plugin::error(ErrorCode::INVALID_ARGUMENT,1005 "cannot find locked buffer");1006 1007 // The address in the entry should be the same we are unregistering.1008 if (Entry->HstPtr != HstPtr)1009 return Plugin::error(ErrorCode::INVALID_ARGUMENT,1010 "unexpected host pointer in locked buffer entry");1011 1012 // Unregister from the entry.1013 auto LastUseOrErr = unregisterEntryUse(*Entry);1014 if (!LastUseOrErr)1015 return LastUseOrErr.takeError();1016 1017 // There should be no other references to the pinned allocation.1018 if (!(*LastUseOrErr))1019 return Plugin::error(ErrorCode::INVALID_ARGUMENT,1020 "the locked buffer is still being used");1021 1022 // Erase the entry from the map.1023 return eraseEntry(*Entry);1024}1025 1026Expected<void *> PinnedAllocationMapTy::lockHostBuffer(void *HstPtr,1027 size_t Size) {1028 assert(HstPtr && "Invalid pointer");1029 assert(Size && "Invalid size");1030 1031 std::lock_guard<std::shared_mutex> Lock(Mutex);1032 1033 const EntryTy *Entry = findIntersecting(HstPtr);1034 1035 if (Entry) {1036 // An already registered intersecting buffer was found. Register a new use.1037 if (auto Err = registerEntryUse(*Entry, HstPtr, Size))1038 return std::move(Err);1039 1040 // Return the device accessible pointer with the correct offset.1041 return utils::advancePtr(Entry->DevAccessiblePtr,1042 utils::getPtrDiff(HstPtr, Entry->HstPtr));1043 }1044 1045 // No intersecting registered allocation found in the map. First, lock the1046 // host buffer and retrieve the device accessible pointer.1047 auto DevAccessiblePtrOrErr = Device.dataLockImpl(HstPtr, Size);1048 if (!DevAccessiblePtrOrErr)1049 return DevAccessiblePtrOrErr.takeError();1050 1051 // Now insert the new entry into the map.1052 if (auto Err = insertEntry(HstPtr, *DevAccessiblePtrOrErr, Size))1053 return std::move(Err);1054 1055 // Return the device accessible pointer.1056 return *DevAccessiblePtrOrErr;1057}1058 1059Error PinnedAllocationMapTy::unlockHostBuffer(void *HstPtr) {1060 assert(HstPtr && "Invalid pointer");1061 1062 std::lock_guard<std::shared_mutex> Lock(Mutex);1063 1064 const EntryTy *Entry = findIntersecting(HstPtr);1065 if (!Entry)1066 return Plugin::error(ErrorCode::INVALID_ARGUMENT,1067 "cannot find locked buffer");1068 1069 // Unregister from the locked buffer. No need to do anything if there are1070 // others using the allocation.1071 auto LastUseOrErr = unregisterEntryUse(*Entry);1072 if (!LastUseOrErr)1073 return LastUseOrErr.takeError();1074 1075 // No need to do anything if there are others using the allocation.1076 if (!(*LastUseOrErr))1077 return Plugin::success();1078 1079 // This was the last user of the allocation. Unlock the original locked buffer1080 // if it was locked by the plugin. Do not unlock it if it was locked by an1081 // external entity. Unlock the buffer using the host pointer of the entry.1082 if (!Entry->ExternallyLocked)1083 if (auto Err = Device.dataUnlockImpl(Entry->HstPtr))1084 return Err;1085 1086 // Erase the entry from the map.1087 return eraseEntry(*Entry);1088}1089 1090Error PinnedAllocationMapTy::lockMappedHostBuffer(void *HstPtr, size_t Size) {1091 assert(HstPtr && "Invalid pointer");1092 assert(Size && "Invalid size");1093 1094 std::lock_guard<std::shared_mutex> Lock(Mutex);1095 1096 // If previously registered, just register a new user on the entry.1097 const EntryTy *Entry = findIntersecting(HstPtr);1098 if (Entry)1099 return registerEntryUse(*Entry, HstPtr, Size);1100 1101 size_t BaseSize;1102 void *BaseHstPtr, *BaseDevAccessiblePtr;1103 1104 // Check if it was externally pinned by a vendor-specific API.1105 auto IsPinnedOrErr = Device.isPinnedPtrImpl(HstPtr, BaseHstPtr,1106 BaseDevAccessiblePtr, BaseSize);1107 if (!IsPinnedOrErr)1108 return IsPinnedOrErr.takeError();1109 1110 // If pinned, just insert the entry representing the whole pinned buffer.1111 if (*IsPinnedOrErr)1112 return insertEntry(BaseHstPtr, BaseDevAccessiblePtr, BaseSize,1113 /*Externallylocked=*/true);1114 1115 // Not externally pinned. Do nothing if locking of mapped buffers is disabled.1116 if (!LockMappedBuffers)1117 return Plugin::success();1118 1119 // Otherwise, lock the buffer and insert the new entry.1120 auto DevAccessiblePtrOrErr = Device.dataLockImpl(HstPtr, Size);1121 if (!DevAccessiblePtrOrErr) {1122 // Errors may be tolerated.1123 if (!IgnoreLockMappedFailures)1124 return DevAccessiblePtrOrErr.takeError();1125 1126 consumeError(DevAccessiblePtrOrErr.takeError());1127 return Plugin::success();1128 }1129 1130 return insertEntry(HstPtr, *DevAccessiblePtrOrErr, Size);1131}1132 1133Error PinnedAllocationMapTy::unlockUnmappedHostBuffer(void *HstPtr) {1134 assert(HstPtr && "Invalid pointer");1135 1136 std::lock_guard<std::shared_mutex> Lock(Mutex);1137 1138 // Check whether there is any intersecting entry.1139 const EntryTy *Entry = findIntersecting(HstPtr);1140 1141 // No entry but automatic locking of mapped buffers is disabled, so1142 // nothing to do.1143 if (!Entry && !LockMappedBuffers)1144 return Plugin::success();1145 1146 // No entry, automatic locking is enabled, but the locking may have failed, so1147 // do nothing.1148 if (!Entry && IgnoreLockMappedFailures)1149 return Plugin::success();1150 1151 // No entry, but the automatic locking is enabled, so this is an error.1152 if (!Entry)1153 return Plugin::error(ErrorCode::INVALID_ARGUMENT,1154 "locked buffer not found");1155 1156 // There is entry, so unregister a user and check whether it was the last one.1157 auto LastUseOrErr = unregisterEntryUse(*Entry);1158 if (!LastUseOrErr)1159 return LastUseOrErr.takeError();1160 1161 // If it is not the last one, there is nothing to do.1162 if (!(*LastUseOrErr))1163 return Plugin::success();1164 1165 // Otherwise, if it was the last and the buffer was locked by the plugin,1166 // unlock it.1167 if (!Entry->ExternallyLocked)1168 if (auto Err = Device.dataUnlockImpl(Entry->HstPtr))1169 return Err;1170 1171 // Finally erase the entry from the map.1172 return eraseEntry(*Entry);1173}1174 1175Error GenericDeviceTy::synchronize(__tgt_async_info *AsyncInfo,1176 bool ReleaseQueue) {1177 if (!AsyncInfo)1178 return Plugin::error(ErrorCode::INVALID_ARGUMENT,1179 "invalid async info queue");1180 1181 SmallVector<void *> AllocsToDelete{};1182 {1183 std::lock_guard<std::mutex> AllocationGuard{AsyncInfo->Mutex};1184 1185 // This can be false when no work has been added to the AsyncInfo. In which1186 // case, the device has nothing to synchronize.1187 if (AsyncInfo->Queue)1188 if (auto Err = synchronizeImpl(*AsyncInfo, ReleaseQueue))1189 return Err;1190 1191 std::swap(AllocsToDelete, AsyncInfo->AssociatedAllocations);1192 }1193 1194 for (auto *Ptr : AllocsToDelete)1195 if (auto Err = dataDelete(Ptr, TargetAllocTy::TARGET_ALLOC_DEVICE))1196 return Err;1197 1198 return Plugin::success();1199}1200 1201Error GenericDeviceTy::queryAsync(__tgt_async_info *AsyncInfo) {1202 if (!AsyncInfo || !AsyncInfo->Queue)1203 return Plugin::error(ErrorCode::INVALID_ARGUMENT,1204 "invalid async info queue");1205 1206 return queryAsyncImpl(*AsyncInfo);1207}1208 1209Error GenericDeviceTy::memoryVAMap(void **Addr, void *VAddr, size_t *RSize) {1210 return Plugin::error(ErrorCode::UNSUPPORTED,1211 "device does not support VA Management");1212}1213 1214Error GenericDeviceTy::memoryVAUnMap(void *VAddr, size_t Size) {1215 return Plugin::error(ErrorCode::UNSUPPORTED,1216 "device does not support VA Management");1217}1218 1219Error GenericDeviceTy::getDeviceMemorySize(uint64_t &DSize) {1220 return Plugin::error(1221 ErrorCode::UNIMPLEMENTED,1222 "missing getDeviceMemorySize implementation (required by RR-heuristic");1223}1224 1225Expected<void *> GenericDeviceTy::dataAlloc(int64_t Size, void *HostPtr,1226 TargetAllocTy Kind) {1227 void *Alloc = nullptr;1228 1229 if (Plugin.getRecordReplay().isRecordingOrReplaying())1230 return Plugin.getRecordReplay().alloc(Size);1231 1232 switch (Kind) {1233 case TARGET_ALLOC_DEFAULT:1234 case TARGET_ALLOC_DEVICE:1235 if (MemoryManager) {1236 auto AllocOrErr = MemoryManager->allocate(Size, HostPtr);1237 if (!AllocOrErr)1238 return AllocOrErr.takeError();1239 Alloc = *AllocOrErr;1240 if (!Alloc)1241 return Plugin::error(ErrorCode::OUT_OF_RESOURCES,1242 "failed to allocate from memory manager");1243 break;1244 }1245 [[fallthrough]];1246 case TARGET_ALLOC_HOST:1247 case TARGET_ALLOC_SHARED: {1248 auto AllocOrErr = allocate(Size, HostPtr, Kind);1249 if (!AllocOrErr)1250 return AllocOrErr.takeError();1251 Alloc = *AllocOrErr;1252 if (!Alloc)1253 return Plugin::error(ErrorCode::OUT_OF_RESOURCES,1254 "failed to allocate from device allocator");1255 }1256 }1257 1258 // Report error if the memory manager or the device allocator did not return1259 // any memory buffer.1260 if (!Alloc)1261 return Plugin::error(ErrorCode::UNIMPLEMENTED,1262 "invalid target data allocation kind or requested "1263 "allocator not implemented yet");1264 1265 // Register allocated buffer as pinned memory if the type is host memory.1266 if (Kind == TARGET_ALLOC_HOST)1267 if (auto Err = PinnedAllocs.registerHostBuffer(Alloc, Alloc, Size))1268 return std::move(Err);1269 1270 // Keep track of the allocation stack if we track allocation traces.1271 if (OMPX_TrackAllocationTraces) {1272 std::string StackTrace;1273 llvm::raw_string_ostream OS(StackTrace);1274 llvm::sys::PrintStackTrace(OS);1275 1276 AllocationTraceInfoTy *ATI = new AllocationTraceInfoTy();1277 ATI->AllocationTrace = std::move(StackTrace);1278 ATI->DevicePtr = Alloc;1279 ATI->HostPtr = HostPtr;1280 ATI->Size = Size;1281 ATI->Kind = Kind;1282 1283 auto AllocationTraceMap = AllocationTraces.getExclusiveAccessor();1284 auto *&MapATI = (*AllocationTraceMap)[Alloc];1285 ATI->LastAllocationInfo = MapATI;1286 MapATI = ATI;1287 }1288 1289 return Alloc;1290}1291 1292Error GenericDeviceTy::dataDelete(void *TgtPtr, TargetAllocTy Kind) {1293 // Free is a noop when recording or replaying.1294 if (Plugin.getRecordReplay().isRecordingOrReplaying())1295 return Plugin::success();1296 1297 // Keep track of the deallocation stack if we track allocation traces.1298 if (OMPX_TrackAllocationTraces) {1299 AllocationTraceInfoTy *ATI = nullptr;1300 {1301 auto AllocationTraceMap = AllocationTraces.getExclusiveAccessor();1302 ATI = (*AllocationTraceMap)[TgtPtr];1303 }1304 1305 std::string StackTrace;1306 llvm::raw_string_ostream OS(StackTrace);1307 llvm::sys::PrintStackTrace(OS);1308 1309 if (!ATI)1310 ErrorReporter::reportDeallocationOfNonAllocatedPtr(TgtPtr, Kind, ATI,1311 StackTrace);1312 1313 // ATI is not null, thus we can lock it to inspect and modify it further.1314 std::lock_guard<std::mutex> LG(ATI->Lock);1315 if (!ATI->DeallocationTrace.empty())1316 ErrorReporter::reportDeallocationOfDeallocatedPtr(TgtPtr, Kind, ATI,1317 StackTrace);1318 1319 if (ATI->Kind != Kind)1320 ErrorReporter::reportDeallocationOfWrongPtrKind(TgtPtr, Kind, ATI,1321 StackTrace);1322 1323 ATI->DeallocationTrace = StackTrace;1324 1325#undef DEALLOCATION_ERROR1326 }1327 1328 switch (Kind) {1329 case TARGET_ALLOC_DEFAULT:1330 case TARGET_ALLOC_DEVICE:1331 if (MemoryManager) {1332 if (auto Err = MemoryManager->free(TgtPtr))1333 return Err;1334 break;1335 }1336 [[fallthrough]];1337 case TARGET_ALLOC_HOST:1338 case TARGET_ALLOC_SHARED:1339 if (auto Err = free(TgtPtr, Kind))1340 return Err;1341 }1342 1343 // Unregister deallocated pinned memory buffer if the type is host memory.1344 if (Kind == TARGET_ALLOC_HOST)1345 if (auto Err = PinnedAllocs.unregisterHostBuffer(TgtPtr))1346 return Err;1347 1348 return Plugin::success();1349}1350 1351Error GenericDeviceTy::dataSubmit(void *TgtPtr, const void *HstPtr,1352 int64_t Size, __tgt_async_info *AsyncInfo) {1353 AsyncInfoWrapperTy AsyncInfoWrapper(*this, AsyncInfo);1354 1355 auto Err = dataSubmitImpl(TgtPtr, HstPtr, Size, AsyncInfoWrapper);1356 AsyncInfoWrapper.finalize(Err);1357 return Err;1358}1359 1360Error GenericDeviceTy::dataRetrieve(void *HstPtr, const void *TgtPtr,1361 int64_t Size, __tgt_async_info *AsyncInfo) {1362 AsyncInfoWrapperTy AsyncInfoWrapper(*this, AsyncInfo);1363 1364 auto Err = dataRetrieveImpl(HstPtr, TgtPtr, Size, AsyncInfoWrapper);1365 AsyncInfoWrapper.finalize(Err);1366 return Err;1367}1368 1369Error GenericDeviceTy::dataExchange(const void *SrcPtr, GenericDeviceTy &DstDev,1370 void *DstPtr, int64_t Size,1371 __tgt_async_info *AsyncInfo) {1372 AsyncInfoWrapperTy AsyncInfoWrapper(*this, AsyncInfo);1373 1374 auto Err = dataExchangeImpl(SrcPtr, DstDev, DstPtr, Size, AsyncInfoWrapper);1375 AsyncInfoWrapper.finalize(Err);1376 return Err;1377}1378 1379Error GenericDeviceTy::dataFill(void *TgtPtr, const void *PatternPtr,1380 int64_t PatternSize, int64_t Size,1381 __tgt_async_info *AsyncInfo) {1382 AsyncInfoWrapperTy AsyncInfoWrapper(*this, AsyncInfo);1383 auto Err =1384 dataFillImpl(TgtPtr, PatternPtr, PatternSize, Size, AsyncInfoWrapper);1385 AsyncInfoWrapper.finalize(Err);1386 return Err;1387}1388 1389Error GenericDeviceTy::launchKernel(void *EntryPtr, void **ArgPtrs,1390 ptrdiff_t *ArgOffsets,1391 KernelArgsTy &KernelArgs,1392 __tgt_async_info *AsyncInfo) {1393 AsyncInfoWrapperTy AsyncInfoWrapper(1394 *this,1395 Plugin.getRecordReplay().isRecordingOrReplaying() ? nullptr : AsyncInfo);1396 1397 GenericKernelTy &GenericKernel =1398 *reinterpret_cast<GenericKernelTy *>(EntryPtr);1399 1400 {1401 std::string StackTrace;1402 if (OMPX_TrackNumKernelLaunches) {1403 llvm::raw_string_ostream OS(StackTrace);1404 llvm::sys::PrintStackTrace(OS);1405 }1406 1407 auto KernelTraceInfoRecord = KernelLaunchTraces.getExclusiveAccessor();1408 (*KernelTraceInfoRecord)1409 .emplace(&GenericKernel, std::move(StackTrace), AsyncInfo);1410 }1411 1412 auto Err = GenericKernel.launch(*this, ArgPtrs, ArgOffsets, KernelArgs,1413 AsyncInfoWrapper);1414 1415 // 'finalize' here to guarantee next record-replay actions are in-sync1416 AsyncInfoWrapper.finalize(Err);1417 1418 RecordReplayTy &RecordReplay = Plugin.getRecordReplay();1419 if (RecordReplay.isRecordingOrReplaying() &&1420 RecordReplay.isSaveOutputEnabled())1421 RecordReplay.saveKernelOutputInfo(GenericKernel.getName());1422 1423 return Err;1424}1425 1426Error GenericDeviceTy::initAsyncInfo(__tgt_async_info **AsyncInfoPtr) {1427 assert(AsyncInfoPtr && "Invalid async info");1428 1429 *AsyncInfoPtr = new __tgt_async_info();1430 1431 AsyncInfoWrapperTy AsyncInfoWrapper(*this, *AsyncInfoPtr);1432 1433 auto Err = initAsyncInfoImpl(AsyncInfoWrapper);1434 AsyncInfoWrapper.finalize(Err);1435 return Err;1436}1437 1438Error GenericDeviceTy::enqueueHostCall(void (*Callback)(void *), void *UserData,1439 __tgt_async_info *AsyncInfo) {1440 AsyncInfoWrapperTy AsyncInfoWrapper(*this, AsyncInfo);1441 1442 auto Err = enqueueHostCallImpl(Callback, UserData, AsyncInfoWrapper);1443 AsyncInfoWrapper.finalize(Err);1444 return Err;1445}1446 1447Expected<InfoTreeNode> GenericDeviceTy::obtainInfo() {1448 auto InfoOrErr = obtainInfoImpl();1449 if (InfoOrErr)1450 InfoOrErr->add("UID", getDeviceUid(), "", DeviceInfo::UID);1451 return InfoOrErr;1452}1453 1454Error GenericDeviceTy::printInfo() {1455 auto InfoOrErr = obtainInfo();1456 1457 // Get the vendor-specific info entries describing the device properties.1458 if (auto Err = InfoOrErr.takeError())1459 return Err;1460 1461 // Print all info entries.1462 InfoOrErr->print();1463 1464 return Plugin::success();1465}1466 1467Error GenericDeviceTy::createEvent(void **EventPtrStorage) {1468 return createEventImpl(EventPtrStorage);1469}1470 1471Error GenericDeviceTy::destroyEvent(void *EventPtr) {1472 return destroyEventImpl(EventPtr);1473}1474 1475Error GenericDeviceTy::recordEvent(void *EventPtr,1476 __tgt_async_info *AsyncInfo) {1477 AsyncInfoWrapperTy AsyncInfoWrapper(*this, AsyncInfo);1478 1479 auto Err = recordEventImpl(EventPtr, AsyncInfoWrapper);1480 AsyncInfoWrapper.finalize(Err);1481 return Err;1482}1483 1484Error GenericDeviceTy::waitEvent(void *EventPtr, __tgt_async_info *AsyncInfo) {1485 AsyncInfoWrapperTy AsyncInfoWrapper(*this, AsyncInfo);1486 1487 auto Err = waitEventImpl(EventPtr, AsyncInfoWrapper);1488 AsyncInfoWrapper.finalize(Err);1489 return Err;1490}1491 1492Expected<bool> GenericDeviceTy::hasPendingWork(__tgt_async_info *AsyncInfo) {1493 AsyncInfoWrapperTy AsyncInfoWrapper(*this, AsyncInfo);1494 auto Res = hasPendingWorkImpl(AsyncInfoWrapper);1495 if (auto Err = Res.takeError()) {1496 AsyncInfoWrapper.finalize(Err);1497 return Err;1498 }1499 1500 auto Err = Plugin::success();1501 AsyncInfoWrapper.finalize(Err);1502 if (Err)1503 return Err;1504 return Res;1505}1506 1507Expected<bool> GenericDeviceTy::isEventComplete(void *Event,1508 __tgt_async_info *AsyncInfo) {1509 AsyncInfoWrapperTy AsyncInfoWrapper(*this, AsyncInfo);1510 auto Res = isEventCompleteImpl(Event, AsyncInfoWrapper);1511 if (auto Err = Res.takeError()) {1512 AsyncInfoWrapper.finalize(Err);1513 return Err;1514 }1515 1516 auto Err = Plugin::success();1517 AsyncInfoWrapper.finalize(Err);1518 if (Err)1519 return Err;1520 return Res;1521}1522 1523Error GenericDeviceTy::syncEvent(void *EventPtr) {1524 return syncEventImpl(EventPtr);1525}1526 1527bool GenericDeviceTy::useAutoZeroCopy() { return useAutoZeroCopyImpl(); }1528 1529Expected<bool> GenericDeviceTy::isAccessiblePtr(const void *Ptr, size_t Size) {1530 return isAccessiblePtrImpl(Ptr, Size);1531}1532 1533void GenericDeviceTy::setDeviceUidFromVendorUid(StringRef VendorUid) {1534 DeviceUid = std::string(Plugin.getName()) + "-" + std::string(VendorUid);1535}1536 1537Error GenericPluginTy::init() {1538 if (Initialized)1539 return Plugin::success();1540 1541 auto NumDevicesOrErr = initImpl();1542 if (!NumDevicesOrErr)1543 return NumDevicesOrErr.takeError();1544 Initialized = true;1545 1546 NumDevices = *NumDevicesOrErr;1547 if (NumDevices == 0)1548 return Plugin::success();1549 1550 assert(Devices.size() == 0 && "Plugin already initialized");1551 Devices.resize(NumDevices, nullptr);1552 1553 GlobalHandler = createGlobalHandler();1554 assert(GlobalHandler && "Invalid global handler");1555 1556 RPCServer = new RPCServerTy(*this);1557 assert(RPCServer && "Invalid RPC server");1558 1559 RecordReplay = new RecordReplayTy();1560 assert(RecordReplay && "Invalid RR interface");1561 1562 return Plugin::success();1563}1564 1565Error GenericPluginTy::deinit() {1566 assert(Initialized && "Plugin was not initialized!");1567 1568 // Deinitialize all active devices.1569 for (int32_t DeviceId = 0; DeviceId < NumDevices; ++DeviceId) {1570 if (Devices[DeviceId]) {1571 if (auto Err = deinitDevice(DeviceId))1572 return Err;1573 }1574 assert(!Devices[DeviceId] && "Device was not deinitialized");1575 }1576 1577 // There is no global handler if no device is available.1578 if (GlobalHandler)1579 delete GlobalHandler;1580 1581 if (RPCServer) {1582 if (Error Err = RPCServer->shutDown())1583 return Err;1584 delete RPCServer;1585 }1586 1587 if (RecordReplay)1588 delete RecordReplay;1589 1590 // Perform last deinitializations on the plugin.1591 if (Error Err = deinitImpl())1592 return Err;1593 Initialized = false;1594 1595 return Plugin::success();1596}1597 1598Error GenericPluginTy::initDevice(int32_t DeviceId) {1599 assert(!Devices[DeviceId] && "Device already initialized");1600 1601 // Create the device and save the reference.1602 GenericDeviceTy *Device = createDevice(*this, DeviceId, NumDevices);1603 assert(Device && "Invalid device");1604 1605 // Save the device reference into the list.1606 Devices[DeviceId] = Device;1607 1608 // Initialize the device and its resources.1609 return Device->init(*this);1610}1611 1612Error GenericPluginTy::deinitDevice(int32_t DeviceId) {1613 // The device may be already deinitialized.1614 if (Devices[DeviceId] == nullptr)1615 return Plugin::success();1616 1617 // Deinitialize the device and release its resources.1618 if (auto Err = Devices[DeviceId]->deinit(*this))1619 return Err;1620 1621 // Delete the device and invalidate its reference.1622 delete Devices[DeviceId];1623 Devices[DeviceId] = nullptr;1624 1625 return Plugin::success();1626}1627 1628Expected<bool> GenericPluginTy::checkELFImage(StringRef Image) const {1629 // First check if this image is a regular ELF file.1630 if (!utils::elf::isELF(Image))1631 return false;1632 1633 // Check if this image is an ELF with a matching machine value.1634 auto MachineOrErr = utils::elf::checkMachine(Image, getMagicElfBits());1635 if (!MachineOrErr)1636 return MachineOrErr.takeError();1637 1638 return MachineOrErr;1639}1640 1641Expected<bool> GenericPluginTy::checkBitcodeImage(StringRef Image) const {1642 if (identify_magic(Image) != file_magic::bitcode)1643 return false;1644 1645 LLVMContext Context;1646 auto ModuleOrErr = getLazyBitcodeModule(MemoryBufferRef(Image, ""), Context,1647 /*ShouldLazyLoadMetadata=*/true);1648 if (!ModuleOrErr)1649 return ModuleOrErr.takeError();1650 Module &M = **ModuleOrErr;1651 1652 return M.getTargetTriple().getArch() == getTripleArch();1653}1654 1655int32_t GenericPluginTy::is_initialized() const { return Initialized; }1656 1657int32_t GenericPluginTy::isPluginCompatible(StringRef Image) {1658 auto HandleError = [&](Error Err) -> bool {1659 [[maybe_unused]] std::string ErrStr = toString(std::move(Err));1660 DP("Failure to check validity of image %p: %s", Image.data(),1661 ErrStr.c_str());1662 return false;1663 };1664 switch (identify_magic(Image)) {1665 case file_magic::elf:1666 case file_magic::elf_relocatable:1667 case file_magic::elf_executable:1668 case file_magic::elf_shared_object:1669 case file_magic::elf_core: {1670 auto MatchOrErr = checkELFImage(Image);1671 if (Error Err = MatchOrErr.takeError())1672 return HandleError(std::move(Err));1673 return *MatchOrErr;1674 }1675 case file_magic::bitcode: {1676 auto MatchOrErr = checkBitcodeImage(Image);1677 if (Error Err = MatchOrErr.takeError())1678 return HandleError(std::move(Err));1679 return *MatchOrErr;1680 }1681 default:1682 return false;1683 }1684}1685 1686int32_t GenericPluginTy::isDeviceCompatible(int32_t DeviceId, StringRef Image) {1687 auto HandleError = [&](Error Err) -> bool {1688 [[maybe_unused]] std::string ErrStr = toString(std::move(Err));1689 DP("Failure to check validity of image %p: %s", Image.data(),1690 ErrStr.c_str());1691 return false;1692 };1693 switch (identify_magic(Image)) {1694 case file_magic::elf:1695 case file_magic::elf_relocatable:1696 case file_magic::elf_executable:1697 case file_magic::elf_shared_object:1698 case file_magic::elf_core: {1699 auto MatchOrErr = checkELFImage(Image);1700 if (Error Err = MatchOrErr.takeError())1701 return HandleError(std::move(Err));1702 if (!*MatchOrErr)1703 return false;1704 1705 // Perform plugin-dependent checks for the specific architecture if needed.1706 auto CompatibleOrErr = isELFCompatible(DeviceId, Image);1707 if (Error Err = CompatibleOrErr.takeError())1708 return HandleError(std::move(Err));1709 return *CompatibleOrErr;1710 }1711 case file_magic::bitcode: {1712 auto MatchOrErr = checkBitcodeImage(Image);1713 if (Error Err = MatchOrErr.takeError())1714 return HandleError(std::move(Err));1715 return *MatchOrErr;1716 }1717 default:1718 return false;1719 }1720}1721 1722int32_t GenericPluginTy::is_device_initialized(int32_t DeviceId) const {1723 return isValidDeviceId(DeviceId) && Devices[DeviceId] != nullptr;1724}1725 1726int32_t GenericPluginTy::init_device(int32_t DeviceId) {1727 auto Err = initDevice(DeviceId);1728 if (Err) {1729 REPORT("Failure to initialize device %d: %s\n", DeviceId,1730 toString(std::move(Err)).data());1731 return OFFLOAD_FAIL;1732 }1733 1734 return OFFLOAD_SUCCESS;1735}1736 1737int32_t GenericPluginTy::number_of_devices() { return getNumDevices(); }1738 1739int32_t GenericPluginTy::is_data_exchangable(int32_t SrcDeviceId,1740 int32_t DstDeviceId) {1741 return isDataExchangable(SrcDeviceId, DstDeviceId);1742}1743 1744int32_t GenericPluginTy::initialize_record_replay(int32_t DeviceId,1745 int64_t MemorySize,1746 void *VAddr, bool isRecord,1747 bool SaveOutput,1748 uint64_t &ReqPtrArgOffset) {1749 GenericDeviceTy &Device = getDevice(DeviceId);1750 RecordReplayTy::RRStatusTy Status =1751 isRecord ? RecordReplayTy::RRStatusTy::RRRecording1752 : RecordReplayTy::RRStatusTy::RRReplaying;1753 1754 if (auto Err = RecordReplay->init(&Device, MemorySize, VAddr, Status,1755 SaveOutput, ReqPtrArgOffset)) {1756 REPORT("WARNING RR did not initialize RR-properly with %lu bytes"1757 "(Error: %s)\n",1758 MemorySize, toString(std::move(Err)).data());1759 RecordReplay->setStatus(RecordReplayTy::RRStatusTy::RRDeactivated);1760 1761 if (!isRecord) {1762 return OFFLOAD_FAIL;1763 }1764 }1765 return OFFLOAD_SUCCESS;1766}1767 1768int32_t GenericPluginTy::load_binary(int32_t DeviceId,1769 __tgt_device_image *TgtImage,1770 __tgt_device_binary *Binary) {1771 GenericDeviceTy &Device = getDevice(DeviceId);1772 1773 StringRef Buffer(reinterpret_cast<const char *>(TgtImage->ImageStart),1774 utils::getPtrDiff(TgtImage->ImageEnd, TgtImage->ImageStart));1775 auto ImageOrErr = Device.loadBinary(*this, Buffer);1776 if (!ImageOrErr) {1777 auto Err = ImageOrErr.takeError();1778 REPORT("Failure to load binary image %p on device %d: %s\n", TgtImage,1779 DeviceId, toString(std::move(Err)).data());1780 return OFFLOAD_FAIL;1781 }1782 1783 DeviceImageTy *Image = *ImageOrErr;1784 assert(Image != nullptr && "Invalid Image");1785 1786 *Binary = __tgt_device_binary{reinterpret_cast<uint64_t>(Image)};1787 1788 return OFFLOAD_SUCCESS;1789}1790 1791void *GenericPluginTy::data_alloc(int32_t DeviceId, int64_t Size, void *HostPtr,1792 int32_t Kind) {1793 auto AllocOrErr =1794 getDevice(DeviceId).dataAlloc(Size, HostPtr, (TargetAllocTy)Kind);1795 if (!AllocOrErr) {1796 auto Err = AllocOrErr.takeError();1797 REPORT("Failure to allocate device memory: %s\n",1798 toString(std::move(Err)).data());1799 return nullptr;1800 }1801 assert(*AllocOrErr && "Null pointer upon successful allocation");1802 1803 return *AllocOrErr;1804}1805 1806int32_t GenericPluginTy::data_delete(int32_t DeviceId, void *TgtPtr,1807 int32_t Kind) {1808 auto Err =1809 getDevice(DeviceId).dataDelete(TgtPtr, static_cast<TargetAllocTy>(Kind));1810 if (Err) {1811 REPORT("Failure to deallocate device pointer %p: %s\n", TgtPtr,1812 toString(std::move(Err)).data());1813 return OFFLOAD_FAIL;1814 }1815 1816 return OFFLOAD_SUCCESS;1817}1818 1819int32_t GenericPluginTy::data_lock(int32_t DeviceId, void *Ptr, int64_t Size,1820 void **LockedPtr) {1821 auto LockedPtrOrErr = getDevice(DeviceId).dataLock(Ptr, Size);1822 if (!LockedPtrOrErr) {1823 auto Err = LockedPtrOrErr.takeError();1824 REPORT("Failure to lock memory %p: %s\n", Ptr,1825 toString(std::move(Err)).data());1826 return OFFLOAD_FAIL;1827 }1828 1829 if (!(*LockedPtrOrErr)) {1830 REPORT("Failure to lock memory %p: obtained a null locked pointer\n", Ptr);1831 return OFFLOAD_FAIL;1832 }1833 *LockedPtr = *LockedPtrOrErr;1834 1835 return OFFLOAD_SUCCESS;1836}1837 1838int32_t GenericPluginTy::data_unlock(int32_t DeviceId, void *Ptr) {1839 auto Err = getDevice(DeviceId).dataUnlock(Ptr);1840 if (Err) {1841 REPORT("Failure to unlock memory %p: %s\n", Ptr,1842 toString(std::move(Err)).data());1843 return OFFLOAD_FAIL;1844 }1845 1846 return OFFLOAD_SUCCESS;1847}1848 1849int32_t GenericPluginTy::data_notify_mapped(int32_t DeviceId, void *HstPtr,1850 int64_t Size) {1851 auto Err = getDevice(DeviceId).notifyDataMapped(HstPtr, Size);1852 if (Err) {1853 REPORT("Failure to notify data mapped %p: %s\n", HstPtr,1854 toString(std::move(Err)).data());1855 return OFFLOAD_FAIL;1856 }1857 1858 return OFFLOAD_SUCCESS;1859}1860 1861int32_t GenericPluginTy::data_notify_unmapped(int32_t DeviceId, void *HstPtr) {1862 auto Err = getDevice(DeviceId).notifyDataUnmapped(HstPtr);1863 if (Err) {1864 REPORT("Failure to notify data unmapped %p: %s\n", HstPtr,1865 toString(std::move(Err)).data());1866 return OFFLOAD_FAIL;1867 }1868 1869 return OFFLOAD_SUCCESS;1870}1871 1872int32_t GenericPluginTy::data_submit(int32_t DeviceId, void *TgtPtr,1873 void *HstPtr, int64_t Size) {1874 return data_submit_async(DeviceId, TgtPtr, HstPtr, Size,1875 /*AsyncInfoPtr=*/nullptr);1876}1877 1878int32_t GenericPluginTy::data_submit_async(int32_t DeviceId, void *TgtPtr,1879 void *HstPtr, int64_t Size,1880 __tgt_async_info *AsyncInfoPtr) {1881 auto Err = getDevice(DeviceId).dataSubmit(TgtPtr, HstPtr, Size, AsyncInfoPtr);1882 if (Err) {1883 REPORT("Failure to copy data from host to device. Pointers: host "1884 "= " DPxMOD ", device = " DPxMOD ", size = %" PRId64 ": %s\n",1885 DPxPTR(HstPtr), DPxPTR(TgtPtr), Size,1886 toString(std::move(Err)).data());1887 return OFFLOAD_FAIL;1888 }1889 1890 return OFFLOAD_SUCCESS;1891}1892 1893int32_t GenericPluginTy::data_retrieve(int32_t DeviceId, void *HstPtr,1894 void *TgtPtr, int64_t Size) {1895 return data_retrieve_async(DeviceId, HstPtr, TgtPtr, Size,1896 /*AsyncInfoPtr=*/nullptr);1897}1898 1899int32_t GenericPluginTy::data_retrieve_async(int32_t DeviceId, void *HstPtr,1900 void *TgtPtr, int64_t Size,1901 __tgt_async_info *AsyncInfoPtr) {1902 auto Err =1903 getDevice(DeviceId).dataRetrieve(HstPtr, TgtPtr, Size, AsyncInfoPtr);1904 if (Err) {1905 REPORT("Failure to copy data from device to host. Pointers: host "1906 "= " DPxMOD ", device = " DPxMOD ", size = %" PRId64 ": %s\n",1907 DPxPTR(HstPtr), DPxPTR(TgtPtr), Size,1908 toString(std::move(Err)).data());1909 return OFFLOAD_FAIL;1910 }1911 1912 return OFFLOAD_SUCCESS;1913}1914 1915int32_t GenericPluginTy::data_exchange(int32_t SrcDeviceId, void *SrcPtr,1916 int32_t DstDeviceId, void *DstPtr,1917 int64_t Size) {1918 return data_exchange_async(SrcDeviceId, SrcPtr, DstDeviceId, DstPtr, Size,1919 /*AsyncInfoPtr=*/nullptr);1920}1921 1922int32_t GenericPluginTy::data_exchange_async(int32_t SrcDeviceId, void *SrcPtr,1923 int DstDeviceId, void *DstPtr,1924 int64_t Size,1925 __tgt_async_info *AsyncInfo) {1926 GenericDeviceTy &SrcDevice = getDevice(SrcDeviceId);1927 GenericDeviceTy &DstDevice = getDevice(DstDeviceId);1928 auto Err = SrcDevice.dataExchange(SrcPtr, DstDevice, DstPtr, Size, AsyncInfo);1929 if (Err) {1930 REPORT("Failure to copy data from device (%d) to device (%d). Pointers: "1931 "host = " DPxMOD ", device = " DPxMOD ", size = %" PRId64 ": %s\n",1932 SrcDeviceId, DstDeviceId, DPxPTR(SrcPtr), DPxPTR(DstPtr), Size,1933 toString(std::move(Err)).data());1934 return OFFLOAD_FAIL;1935 }1936 1937 return OFFLOAD_SUCCESS;1938}1939 1940int32_t GenericPluginTy::launch_kernel(int32_t DeviceId, void *TgtEntryPtr,1941 void **TgtArgs, ptrdiff_t *TgtOffsets,1942 KernelArgsTy *KernelArgs,1943 __tgt_async_info *AsyncInfoPtr) {1944 auto Err = getDevice(DeviceId).launchKernel(TgtEntryPtr, TgtArgs, TgtOffsets,1945 *KernelArgs, AsyncInfoPtr);1946 if (Err) {1947 REPORT("Failure to run target region " DPxMOD " in device %d: %s\n",1948 DPxPTR(TgtEntryPtr), DeviceId, toString(std::move(Err)).data());1949 return OFFLOAD_FAIL;1950 }1951 1952 return OFFLOAD_SUCCESS;1953}1954 1955int32_t GenericPluginTy::synchronize(int32_t DeviceId,1956 __tgt_async_info *AsyncInfoPtr) {1957 auto Err = getDevice(DeviceId).synchronize(AsyncInfoPtr);1958 if (Err) {1959 REPORT("Failure to synchronize stream %p: %s\n", AsyncInfoPtr->Queue,1960 toString(std::move(Err)).data());1961 return OFFLOAD_FAIL;1962 }1963 1964 return OFFLOAD_SUCCESS;1965}1966 1967int32_t GenericPluginTy::query_async(int32_t DeviceId,1968 __tgt_async_info *AsyncInfoPtr) {1969 auto Err = getDevice(DeviceId).queryAsync(AsyncInfoPtr);1970 if (Err) {1971 REPORT("Failure to query stream %p: %s\n", AsyncInfoPtr->Queue,1972 toString(std::move(Err)).data());1973 return OFFLOAD_FAIL;1974 }1975 1976 return OFFLOAD_SUCCESS;1977}1978 1979void GenericPluginTy::print_device_info(int32_t DeviceId) {1980 if (auto Err = getDevice(DeviceId).printInfo())1981 REPORT("Failure to print device %d info: %s\n", DeviceId,1982 toString(std::move(Err)).data());1983}1984 1985int32_t GenericPluginTy::create_event(int32_t DeviceId, void **EventPtr) {1986 auto Err = getDevice(DeviceId).createEvent(EventPtr);1987 if (Err) {1988 REPORT("Failure to create event: %s\n", toString(std::move(Err)).data());1989 return OFFLOAD_FAIL;1990 }1991 1992 return OFFLOAD_SUCCESS;1993}1994 1995int32_t GenericPluginTy::record_event(int32_t DeviceId, void *EventPtr,1996 __tgt_async_info *AsyncInfoPtr) {1997 auto Err = getDevice(DeviceId).recordEvent(EventPtr, AsyncInfoPtr);1998 if (Err) {1999 REPORT("Failure to record event %p: %s\n", EventPtr,2000 toString(std::move(Err)).data());2001 return OFFLOAD_FAIL;2002 }2003 2004 return OFFLOAD_SUCCESS;2005}2006 2007int32_t GenericPluginTy::wait_event(int32_t DeviceId, void *EventPtr,2008 __tgt_async_info *AsyncInfoPtr) {2009 auto Err = getDevice(DeviceId).waitEvent(EventPtr, AsyncInfoPtr);2010 if (Err) {2011 REPORT("Failure to wait event %p: %s\n", EventPtr,2012 toString(std::move(Err)).data());2013 return OFFLOAD_FAIL;2014 }2015 2016 return OFFLOAD_SUCCESS;2017}2018 2019int32_t GenericPluginTy::sync_event(int32_t DeviceId, void *EventPtr) {2020 auto Err = getDevice(DeviceId).syncEvent(EventPtr);2021 if (Err) {2022 REPORT("Failure to synchronize event %p: %s\n", EventPtr,2023 toString(std::move(Err)).data());2024 return OFFLOAD_FAIL;2025 }2026 2027 return OFFLOAD_SUCCESS;2028}2029 2030int32_t GenericPluginTy::destroy_event(int32_t DeviceId, void *EventPtr) {2031 auto Err = getDevice(DeviceId).destroyEvent(EventPtr);2032 if (Err) {2033 REPORT("Failure to destroy event %p: %s\n", EventPtr,2034 toString(std::move(Err)).data());2035 return OFFLOAD_FAIL;2036 }2037 2038 return OFFLOAD_SUCCESS;2039}2040 2041void GenericPluginTy::set_info_flag(uint32_t NewInfoLevel) {2042 std::atomic<uint32_t> &InfoLevel = getInfoLevelInternal();2043 InfoLevel.store(NewInfoLevel);2044}2045 2046int32_t GenericPluginTy::init_async_info(int32_t DeviceId,2047 __tgt_async_info **AsyncInfoPtr) {2048 assert(AsyncInfoPtr && "Invalid async info");2049 2050 auto Err = getDevice(DeviceId).initAsyncInfo(AsyncInfoPtr);2051 if (Err) {2052 REPORT("Failure to initialize async info at " DPxMOD " on device %d: %s\n",2053 DPxPTR(*AsyncInfoPtr), DeviceId, toString(std::move(Err)).data());2054 return OFFLOAD_FAIL;2055 }2056 2057 return OFFLOAD_SUCCESS;2058}2059 2060int32_t GenericPluginTy::set_device_identifier(int32_t UserId,2061 int32_t DeviceId) {2062 UserDeviceIds[DeviceId] = UserId;2063 2064 return OFFLOAD_SUCCESS;2065}2066 2067int32_t GenericPluginTy::use_auto_zero_copy(int32_t DeviceId) {2068 return getDevice(DeviceId).useAutoZeroCopy();2069}2070 2071int32_t GenericPluginTy::is_accessible_ptr(int32_t DeviceId, const void *Ptr,2072 size_t Size) {2073 auto HandleError = [&](Error Err) -> bool {2074 [[maybe_unused]] std::string ErrStr = toString(std::move(Err));2075 DP("Failure while checking accessibility of pointer %p for device %d: %s",2076 Ptr, DeviceId, ErrStr.c_str());2077 return false;2078 };2079 2080 auto AccessibleOrErr = getDevice(DeviceId).isAccessiblePtr(Ptr, Size);2081 if (Error Err = AccessibleOrErr.takeError())2082 return HandleError(std::move(Err));2083 2084 return *AccessibleOrErr;2085}2086 2087int32_t GenericPluginTy::get_global(__tgt_device_binary Binary, uint64_t Size,2088 const char *Name, void **DevicePtr) {2089 assert(Binary.handle && "Invalid device binary handle");2090 DeviceImageTy &Image = *reinterpret_cast<DeviceImageTy *>(Binary.handle);2091 2092 GenericDeviceTy &Device = Image.getDevice();2093 2094 GlobalTy DeviceGlobal(Name, Size);2095 GenericGlobalHandlerTy &GHandler = getGlobalHandler();2096 if (auto Err =2097 GHandler.getGlobalMetadataFromDevice(Device, Image, DeviceGlobal)) {2098 consumeError(std::move(Err));2099 return OFFLOAD_FAIL;2100 }2101 2102 *DevicePtr = DeviceGlobal.getPtr();2103 assert(DevicePtr && "Invalid device global's address");2104 2105 // Save the loaded globals if we are recording.2106 RecordReplayTy &RecordReplay = Device.Plugin.getRecordReplay();2107 if (RecordReplay.isRecording())2108 RecordReplay.addEntry(Name, Size, *DevicePtr);2109 2110 return OFFLOAD_SUCCESS;2111}2112 2113int32_t GenericPluginTy::get_function(__tgt_device_binary Binary,2114 const char *Name, void **KernelPtr) {2115 assert(Binary.handle && "Invalid device binary handle");2116 DeviceImageTy &Image = *reinterpret_cast<DeviceImageTy *>(Binary.handle);2117 2118 GenericDeviceTy &Device = Image.getDevice();2119 2120 auto KernelOrErr = Device.constructKernel(Name);2121 if (Error Err = KernelOrErr.takeError()) {2122 REPORT("Failure to look up kernel: %s\n", toString(std::move(Err)).data());2123 return OFFLOAD_FAIL;2124 }2125 2126 GenericKernelTy &Kernel = *KernelOrErr;2127 if (auto Err = Kernel.init(Device, Image)) {2128 REPORT("Failure to init kernel: %s\n", toString(std::move(Err)).data());2129 return OFFLOAD_FAIL;2130 }2131 2132 // Note that this is not the kernel's device address.2133 *KernelPtr = &Kernel;2134 return OFFLOAD_SUCCESS;2135}2136 2137/// Create OpenMP interop with the given interop context2138omp_interop_val_t *2139GenericPluginTy::create_interop(int32_t ID, int32_t InteropContext,2140 interop_spec_t *InteropSpec) {2141 assert(InteropSpec && "Interop spec is null");2142 auto &Device = getDevice(ID);2143 auto InteropOrErr = Device.createInterop(InteropContext, *InteropSpec);2144 if (!InteropOrErr) {2145 REPORT("Failure to create interop object for device " DPxMOD ": %s\n",2146 DPxPTR(InteropSpec), toString(InteropOrErr.takeError()).c_str());2147 return nullptr;2148 }2149 return *InteropOrErr;2150}2151 2152/// Release OpenMP interop object2153int32_t GenericPluginTy::release_interop(int32_t ID,2154 omp_interop_val_t *Interop) {2155 assert(Interop && "Interop is null");2156 assert(Interop->device_id == ID && "Interop does not match device id");2157 auto &Device = getDevice(ID);2158 auto Err = Device.releaseInterop(Interop);2159 if (Err) {2160 REPORT("Failure to release interop object " DPxMOD ": %s\n",2161 DPxPTR(Interop), toString(std::move(Err)).c_str());2162 return OFFLOAD_FAIL;2163 }2164 return OFFLOAD_SUCCESS;2165}2166 2167/// Flush the queue associated with the interop object if necessary2168int32_t GenericPluginTy::flush_queue(omp_interop_val_t *Interop) {2169 assert(Interop && "Interop is null");2170 auto Err = flushQueueImpl(Interop);2171 if (Err) {2172 REPORT("Failure to flush interop object " DPxMOD " queue: %s\n",2173 DPxPTR(Interop), toString(std::move(Err)).c_str());2174 return OFFLOAD_FAIL;2175 }2176 return OFFLOAD_SUCCESS;2177}2178 2179/// Perform a host synchronization with the queue associated with the interop2180/// object and wait for it to complete.2181int32_t GenericPluginTy::sync_barrier(omp_interop_val_t *Interop) {2182 assert(Interop && "Interop is null");2183 auto Err = syncBarrierImpl(Interop);2184 if (Err) {2185 REPORT("Failure to synchronize interop object " DPxMOD ": %s\n",2186 DPxPTR(Interop), toString(std::move(Err)).c_str());2187 return OFFLOAD_FAIL;2188 }2189 return OFFLOAD_SUCCESS;2190}2191 2192/// Queue an asynchronous barrier in the queue associated with the interop2193/// object and return immediately.2194int32_t GenericPluginTy::async_barrier(omp_interop_val_t *Interop) {2195 assert(Interop && "Interop is null");2196 auto Err = asyncBarrierImpl(Interop);2197 if (Err) {2198 REPORT("Failure to queue barrier in interop object " DPxMOD ": %s\n",2199 DPxPTR(Interop), toString(std::move(Err)).c_str());2200 return OFFLOAD_FAIL;2201 }2202 return OFFLOAD_SUCCESS;2203}2204 2205int32_t GenericPluginTy::data_fence(int32_t DeviceId,2206 __tgt_async_info *AsyncInfo) {2207 auto Err = getDevice(DeviceId).dataFence(AsyncInfo);2208 if (Err) {2209 REPORT("failure to place data fence on device %d: %s\n", DeviceId,2210 toString(std::move(Err)).data());2211 return OFFLOAD_FAIL;2212 }2213 2214 return OFFLOAD_SUCCESS;2215}2216