| //===- PluginInterface.cpp - Target independent plugin device interface ---===// | |
| // | |
| // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. | |
| // See https://llvm.org/LICENSE.txt for license information. | |
| // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception | |
| // | |
| //===----------------------------------------------------------------------===// | |
| // | |
| //===----------------------------------------------------------------------===// | |
| #include "PluginInterface.h" | |
| #include "Shared/APITypes.h" | |
| #include "Shared/Debug.h" | |
| #include "Shared/Environment.h" | |
| #include "ErrorReporting.h" | |
| #include "GlobalHandler.h" | |
| #include "JIT.h" | |
| #include "Shared/Utils.h" | |
| #include "Utils/ELF.h" | |
| #include "omptarget.h" | |
| #ifdef OMPT_SUPPORT | |
| #include "OpenMP/OMPT/Callback.h" | |
| #include "omp-tools.h" | |
| #endif | |
| #include "llvm/Bitcode/BitcodeReader.h" | |
| #include "llvm/Frontend/OpenMP/OMPConstants.h" | |
| #include "llvm/Support/Error.h" | |
| #include "llvm/Support/MathExtras.h" | |
| #include "llvm/Support/MemoryBuffer.h" | |
| #include "llvm/Support/Signals.h" | |
| #include "llvm/Support/raw_ostream.h" | |
| #include <cstdint> | |
| #include <limits> | |
| using namespace llvm; | |
| using namespace omp; | |
| using namespace target; | |
| using namespace plugin; | |
| using namespace error; | |
| using namespace llvm::offload::debug; | |
| AsyncInfoWrapperTy::AsyncInfoWrapperTy(GenericDeviceTy &Device, | |
| __tgt_async_info *AsyncInfoPtr) | |
| : Device(Device), | |
| AsyncInfoPtr(AsyncInfoPtr ? AsyncInfoPtr : &LocalAsyncInfo) {} | |
| Error AsyncInfoWrapperTy::synchronize() { | |
| assert(AsyncInfoPtr && "AsyncInfoWrapperTy already finalized"); | |
| // Synchronize with the async info's operations without releasing the queue. | |
| return Device.synchronize(AsyncInfoPtr, /*ReleaseQueue=*/false); | |
| } | |
| void AsyncInfoWrapperTy::finalize(Error &Err) { | |
| assert(AsyncInfoPtr && "AsyncInfoWrapperTy already finalized"); | |
| // If we used a local async info object we want synchronous behavior. In that | |
| // case, and assuming the current status code is correct, we will synchronize | |
| // explicitly when the object is deleted. Update the error with the result of | |
| // the synchronize operation. | |
| if (AsyncInfoPtr == &LocalAsyncInfo && LocalAsyncInfo.Queue && !Err) | |
| Err = Device.synchronize(&LocalAsyncInfo); | |
| // Invalidate the wrapper object. | |
| AsyncInfoPtr = nullptr; | |
| } | |
| Error GenericKernelTy::init(GenericDeviceTy &GenericDevice, | |
| DeviceImageTy &Image) { | |
| ImagePtr = &Image; | |
| // Retrieve kernel environment object for the kernel. | |
| std::string EnvironmentName = std::string(Name) + "_kernel_environment"; | |
| GenericGlobalHandlerTy &GHandler = GenericDevice.Plugin.getGlobalHandler(); | |
| if (GHandler.isSymbolInImage(GenericDevice, Image, EnvironmentName)) { | |
| GlobalTy KernelEnv(EnvironmentName, sizeof(KernelEnvironment), | |
| &KernelEnvironment); | |
| if (auto Err = | |
| GHandler.readGlobalFromImage(GenericDevice, *ImagePtr, KernelEnv)) | |
| return Err; | |
| } else { | |
| KernelEnvironment = KernelEnvironmentTy{}; | |
| ODBG(OLDT_Kernel) << "Failed to read kernel environment for '" << getName() | |
| << "' Using default Bare (0) execution mode"; | |
| } | |
| // Max = Config.Max > 0 ? min(Config.Max, Device.Max) : Device.Max; | |
| MaxNumThreads = KernelEnvironment.Configuration.MaxThreads > 0 | |
| ? std::min(KernelEnvironment.Configuration.MaxThreads, | |
| int32_t(GenericDevice.getThreadLimit())) | |
| : GenericDevice.getThreadLimit(); | |
| // Pref = Config.Pref > 0 ? max(Config.Pref, Device.Pref) : Device.Pref; | |
| PreferredNumThreads = | |
| KernelEnvironment.Configuration.MinThreads > 0 | |
| ? std::max(KernelEnvironment.Configuration.MinThreads, | |
| int32_t(GenericDevice.getDefaultNumThreads())) | |
| : GenericDevice.getDefaultNumThreads(); | |
| return initImpl(GenericDevice, Image); | |
| } | |
| Expected<KernelLaunchEnvironmentTy *> | |
| GenericKernelTy::getKernelLaunchEnvironment( | |
| GenericDeviceTy &GenericDevice, const KernelArgsTy &KernelArgs, | |
| const DynBlockMemConfTy &DynBlockMemConf, | |
| AsyncInfoWrapperTy &AsyncInfoWrapper, uint32_t NumBlocks0) const { | |
| // Ctor/Dtor have no arguments, replaying uses the original kernel launch | |
| // environment. Older versions of the compiler do not generate a kernel | |
| // launch environment. | |
| if ((GenericDevice.getRecordReplay() && | |
| GenericDevice.getRecordReplay()->isReplaying()) || | |
| KernelArgs.Version < OMP_KERNEL_ARG_MIN_VERSION_WITH_DYN_PTR) | |
| return nullptr; | |
| const auto &RedCfg = KernelEnvironment.Configuration; | |
| const bool NeedsReductionBuffer = RedCfg.ReductionDataSize != 0; | |
| if (NeedsReductionBuffer && KernelArgs.Version < OMP_KERNEL_ARG_VERSION) | |
| return Plugin::error(ErrorCode::INVALID_BINARY, | |
| "kernel was built against an older OpenMP " | |
| "kernel-launch-environment ABI (v%u); current " | |
| "runtime requires v%u for cross-team reductions", | |
| KernelArgs.Version, OMP_KERNEL_ARG_VERSION); | |
| if (!NeedsReductionBuffer && !KernelArgs.DynCGroupMem) | |
| return reinterpret_cast<KernelLaunchEnvironmentTy *>(~0); | |
| auto AllocOrErr = GenericDevice.dataAlloc( | |
| sizeof(KernelLaunchEnvironmentTy), | |
| /*HostPtr=*/nullptr, TargetAllocTy::TARGET_ALLOC_DEVICE, /*Alignment=*/0); | |
| if (!AllocOrErr) | |
| return AllocOrErr.takeError(); | |
| // Remember to free the memory later. | |
| AsyncInfoWrapper.freeAllocationAfterSynchronization(*AllocOrErr); | |
| /// Use the KLE in the __tgt_async_info to ensure a stable address for the | |
| /// async data transfer. | |
| auto &LocalKLE = (*AsyncInfoWrapper).KernelLaunchEnvironment; | |
| LocalKLE = KernelLaunchEnvironment; | |
| LocalKLE.DynCGroupMemSize = DynBlockMemConf.Size; | |
| LocalKLE.DynCGroupMemFbPtr = DynBlockMemConf.FallbackPtr; | |
| LocalKLE.DynCGroupMemFb = DynBlockMemConf.Fallback; | |
| LocalKLE.ReductionBuffer = nullptr; | |
| if (NeedsReductionBuffer) { | |
| // Use number of teams many buffer elements. | |
| auto AllocOrErr = GenericDevice.dataAlloc( | |
| uint64_t(RedCfg.ReductionDataSize) * NumBlocks0, | |
| /*HostPtr=*/nullptr, TargetAllocTy::TARGET_ALLOC_DEVICE, | |
| /*Alignment=*/0); | |
| if (!AllocOrErr) | |
| return AllocOrErr.takeError(); | |
| LocalKLE.ReductionBuffer = *AllocOrErr; | |
| // Remember to free the memory later. | |
| AsyncInfoWrapper.freeAllocationAfterSynchronization(*AllocOrErr); | |
| } | |
| INFO(OMP_INFOTYPE_DATA_TRANSFER, GenericDevice.getDeviceId(), | |
| "Copying data from host to device, HstPtr=" DPxMOD ", TgtPtr=" DPxMOD | |
| ", Size=%" PRId64 ", Name=KernelLaunchEnv\n", | |
| DPxPTR(&LocalKLE), DPxPTR(*AllocOrErr), | |
| sizeof(KernelLaunchEnvironmentTy)); | |
| auto Err = GenericDevice.dataSubmit(*AllocOrErr, &LocalKLE, | |
| sizeof(KernelLaunchEnvironmentTy), | |
| AsyncInfoWrapper); | |
| if (Err) | |
| return Err; | |
| return static_cast<KernelLaunchEnvironmentTy *>(*AllocOrErr); | |
| } | |
| Error GenericKernelTy::printLaunchInfo(GenericDeviceTy &GenericDevice, | |
| KernelArgsTy &KernelArgs, | |
| uint32_t NumThreads[3], | |
| uint32_t NumBlocks[3]) const { | |
| INFO(OMP_INFOTYPE_PLUGIN_KERNEL, GenericDevice.getDeviceId(), | |
| "Launching kernel %s with [%u,%u,%u] blocks and [%u,%u,%u] threads in " | |
| "%s mode\n", | |
| getName(), NumBlocks[0], NumBlocks[1], NumBlocks[2], NumThreads[0], | |
| NumThreads[1], NumThreads[2], getExecutionModeName()); | |
| return printLaunchInfoDetails(GenericDevice, KernelArgs, NumThreads, | |
| NumBlocks); | |
| } | |
| Error GenericKernelTy::printLaunchInfoDetails(GenericDeviceTy &GenericDevice, | |
| KernelArgsTy &KernelArgs, | |
| uint32_t NumThreads[3], | |
| uint32_t NumBlocks[3]) const { | |
| return Plugin::success(); | |
| } | |
| Expected<DynBlockMemConfTy> | |
| GenericKernelTy::prepareBlockMemory(GenericDeviceTy &GenericDevice, | |
| KernelArgsTy &KernelArgs, | |
| uint32_t NumBlocks) const { | |
| uint32_t MaxBlockMemSize = GenericDevice.getMaxBlockSharedMemSize(); | |
| uint32_t DynBlockMemSize = KernelArgs.DynCGroupMem; | |
| uint32_t TotalBlockMemSize = StaticBlockMemSize + DynBlockMemSize; | |
| uint32_t DynNativeBlockMemSize = DynBlockMemSize; | |
| void *DynFallbackPtr = nullptr; | |
| // No enough block memory to cover the static one. Cannot run the kernel. | |
| if (StaticBlockMemSize > MaxBlockMemSize) | |
| return Plugin::error(ErrorCode::INVALID_ARGUMENT, | |
| "Static block memory size exceeds maximum"); | |
| // No enough block memory to cover dynamic one, and the fallback is aborting. | |
| if (static_cast<DynCGroupMemFallbackType>( | |
| KernelArgs.Flags.DynCGroupMemFallback) == | |
| DynCGroupMemFallbackType::Abort && | |
| TotalBlockMemSize > MaxBlockMemSize) | |
| return Plugin::error( | |
| ErrorCode::INVALID_ARGUMENT, | |
| "Requested block memory size (static + dynamic) exceeds maximum"); | |
| DynCGroupMemFallbackType DynFallback = DynCGroupMemFallbackType::None; | |
| if (DynBlockMemSize && TotalBlockMemSize > MaxBlockMemSize) { | |
| // Launch without native dynamic block memory. | |
| DynNativeBlockMemSize = 0; | |
| DynFallback = static_cast<DynCGroupMemFallbackType>( | |
| KernelArgs.Flags.DynCGroupMemFallback); | |
| if (DynFallback != DynCGroupMemFallbackType::DefaultMem) { | |
| // Do not provide any memory as fallback. | |
| DynBlockMemSize = 0; | |
| } else { | |
| // Get global memory as fallback. | |
| auto AllocOrErr = GenericDevice.dataAlloc( | |
| NumBlocks * DynBlockMemSize, | |
| /*HostPtr=*/nullptr, TargetAllocTy::TARGET_ALLOC_DEVICE, | |
| /*Alignment=*/0); | |
| if (!AllocOrErr) | |
| return AllocOrErr.takeError(); | |
| DynFallbackPtr = *AllocOrErr; | |
| } | |
| } | |
| return DynBlockMemConfTy{DynBlockMemSize, DynNativeBlockMemSize, DynFallback, | |
| DynFallbackPtr}; | |
| } | |
| Error GenericKernelTy::launch(GenericDeviceTy &GenericDevice, void **ArgPtrs, | |
| ptrdiff_t *ArgOffsets, KernelArgsTy &KernelArgs, | |
| KernelExtraArgsTy *KernelExtraArgs, | |
| AsyncInfoWrapperTy &AsyncInfoWrapper) const { | |
| llvm::SmallVector<void *, 16> Args; | |
| llvm::SmallVector<void *, 16> Ptrs; | |
| uint32_t EffectiveNumThreads[3] = {KernelArgs.UserThreadLimit[0], | |
| KernelArgs.UserThreadLimit[1], | |
| KernelArgs.UserThreadLimit[2]}; | |
| uint32_t EffectiveNumBlocks[3] = {KernelArgs.UserNumBlocks[0], | |
| KernelArgs.UserNumBlocks[1], | |
| KernelArgs.UserNumBlocks[2]}; | |
| // Multidimensional is only supported with bare mode for now. | |
| assert(isBareMode() || | |
| EffectiveNumThreads[1] == 1 && EffectiveNumThreads[2] == 1 && | |
| EffectiveNumBlocks[1] == 1 && EffectiveNumBlocks[2] == 1 && | |
| "Non-bare mode should only use the first thread and block " | |
| "dimensions"); | |
| assert(!KernelArgs.Flags.StrictBlocksAndThreads || | |
| EffectiveNumThreads[0] > 0 && EffectiveNumThreads[1] > 0 && | |
| EffectiveNumThreads[2] > 0 && EffectiveNumBlocks[0] > 0 && | |
| EffectiveNumBlocks[1] > 0 && EffectiveNumBlocks[2] > 0 && | |
| "Strict requires number of blocks and threads greater than zero"); | |
| // Calculate or adjust the effective number of threads and blocks if needed. | |
| if (!KernelArgs.Flags.StrictBlocksAndThreads) { | |
| EffectiveNumThreads[0] = | |
| getEffectiveNumThreads(GenericDevice, EffectiveNumThreads[0]); | |
| EffectiveNumBlocks[0] = getEffectiveNumBlocks( | |
| GenericDevice, EffectiveNumBlocks[0], KernelArgs.Tripcount, | |
| EffectiveNumThreads[0], KernelArgs.UserThreadLimit[0] > 0); | |
| } | |
| auto DynBlockMemConfOrErr = prepareBlockMemory( | |
| GenericDevice, KernelArgs, | |
| EffectiveNumBlocks[0] * EffectiveNumBlocks[1] * EffectiveNumBlocks[2]); | |
| if (!DynBlockMemConfOrErr) | |
| return DynBlockMemConfOrErr.takeError(); | |
| DynBlockMemConfTy &DynBlockMemConf = *DynBlockMemConfOrErr; | |
| if (DynBlockMemConf.FallbackPtr) | |
| AsyncInfoWrapper.freeAllocationAfterSynchronization( | |
| DynBlockMemConf.FallbackPtr); | |
| auto KernelLaunchEnvOrErr = | |
| getKernelLaunchEnvironment(GenericDevice, KernelArgs, DynBlockMemConf, | |
| AsyncInfoWrapper, EffectiveNumBlocks[0]); | |
| if (!KernelLaunchEnvOrErr) | |
| return KernelLaunchEnvOrErr.takeError(); | |
| KernelLaunchParamsTy LaunchParams; | |
| // Kernel languages do not use the OpenMP indirection and argument parsing. | |
| if (KernelArgs.Flags.IsCUDA) { | |
| LaunchParams = | |
| *reinterpret_cast<KernelLaunchParamsTy *>(KernelArgs.ArgPtrs); | |
| } else if (KernelArgs.Flags.IsPtrArgs) { | |
| LaunchParams = KernelLaunchParamsTy{KernelArgs.NumArgs, KernelArgs.ArgPtrs}; | |
| } else { | |
| LaunchParams = | |
| prepareArgs(GenericDevice, ArgPtrs, ArgOffsets, KernelArgs.NumArgs, | |
| Args, Ptrs, *KernelLaunchEnvOrErr, KernelArgs.Version); | |
| } | |
| if (auto Err = printLaunchInfo(GenericDevice, KernelArgs, EffectiveNumThreads, | |
| EffectiveNumBlocks)) | |
| return Err; | |
| RecordReplayTy::HandleTy RRHandle; | |
| RecordReplayTy *RecordReplay = GenericDevice.getRecordReplay(); | |
| if (RecordReplay) { | |
| // Record replay requires synchronization of any previous operation. | |
| if (auto Err = AsyncInfoWrapper.synchronize()) | |
| return Err; | |
| // Record the kernel prologue data before kernel launch. | |
| auto RRHandleOrErr = RecordReplay->recordPrologue( | |
| *this, KernelArgs, KernelExtraArgs, LaunchParams, EffectiveNumBlocks, | |
| EffectiveNumThreads, DynBlockMemConf.NativeSize); | |
| if (!RRHandleOrErr) | |
| return RRHandleOrErr.takeError(); | |
| RRHandle = *RRHandleOrErr; | |
| } | |
| if (auto Err = launchImpl(GenericDevice, EffectiveNumThreads, | |
| EffectiveNumBlocks, DynBlockMemConf.NativeSize, | |
| KernelArgs, LaunchParams, AsyncInfoWrapper)) | |
| return Err; | |
| if (RecordReplay) { | |
| // Record replay requires synchronization. | |
| if (auto Err = AsyncInfoWrapper.synchronize()) | |
| return Err; | |
| // Record the epilogue data after kernel synchronization. | |
| return RecordReplay->recordEpilogue(*this, RRHandle); | |
| } | |
| return Plugin::success(); | |
| } | |
| KernelLaunchParamsTy | |
| GenericKernelTy::prepareArgs(GenericDeviceTy &GenericDevice, void **ArgPtrs, | |
| ptrdiff_t *ArgOffsets, uint32_t &NumArgs, | |
| llvm::SmallVectorImpl<void *> &Args, | |
| llvm::SmallVectorImpl<void *> &Ptrs, | |
| KernelLaunchEnvironmentTy *KernelLaunchEnvironment, | |
| uint32_t Version) const { | |
| if (NumArgs == 0) | |
| return KernelLaunchParamsTy{}; | |
| // The argument arrays already include the dyn_ptr slot at the end (appended | |
| // by the host for version >= 4, or by upgradeKernelArgs for version 3). | |
| Args.resize(NumArgs); | |
| Ptrs.resize(NumArgs); | |
| for (uint32_t I = 0; I < NumArgs; ++I) | |
| Args[I] = reinterpret_cast<void *>(reinterpret_cast<intptr_t>(ArgPtrs[I]) + | |
| ArgOffsets[I]); | |
| // Optionally assign the KernelLaunchEnvironment to the last slot (dyn_ptr). | |
| if (KernelLaunchEnvironment) | |
| Args[NumArgs - 1] = KernelLaunchEnvironment; | |
| // Version 3 device kernels have dyn_ptr baked in at position 0. Rotate the | |
| // last element to the front to match the device ABI. | |
| if (Version <= OMP_KERNEL_ARG_MIN_VERSION_WITH_DYN_PTR && | |
| KernelLaunchEnvironment) | |
| std::rotate(Args.begin(), Args.end() - 1, Args.end()); | |
| for (uint32_t I = 0; I < NumArgs; ++I) | |
| Ptrs[I] = &Args[I]; | |
| return KernelLaunchParamsTy{NumArgs, &Ptrs[0]}; | |
| } | |
| uint32_t | |
| GenericKernelTy::getEffectiveNumThreads(GenericDeviceTy &GenericDevice, | |
| uint32_t UserThreadLimit) const { | |
| assert(!isBareMode() && "bare kernel should not call this function"); | |
| if (UserThreadLimit > 0 && isGenericMode()) | |
| UserThreadLimit += GenericDevice.getWarpSize(); | |
| return std::min(MaxNumThreads, (UserThreadLimit > 0) ? UserThreadLimit | |
| : PreferredNumThreads); | |
| } | |
| uint32_t GenericKernelTy::getEffectiveNumBlocks( | |
| GenericDeviceTy &GenericDevice, uint32_t UserNumBlocks, | |
| uint64_t LoopTripCount, uint32_t &EffectiveNumThreads, | |
| bool IsNumThreadsFromUser) const { | |
| assert(!isBareMode() && "bare kernel should not call this function"); | |
| // NOTE: This clamps the user-requested number of blocks to the device limit | |
| // rather than honoring it exactly, which is non-standard behavior. Truly | |
| // honoring an arbitrary value would require launching multiple kernels or | |
| // reusing blocks until the requested count has been served. | |
| if (UserNumBlocks > 0) | |
| return std::min(UserNumBlocks, | |
| GenericDevice.getBlockLimit(EffectiveNumThreads)); | |
| // Return the number of blocks required to cover the loop iterations. | |
| if (isNoLoopMode()) | |
| return LoopTripCount > 0 ? (((LoopTripCount - 1) / EffectiveNumThreads) + 1) | |
| : 1; | |
| uint64_t DefaultNumBlocks = GenericDevice.getDefaultNumBlocks(); | |
| uint64_t TripCountNumBlocks = std::numeric_limits<uint64_t>::max(); | |
| if (LoopTripCount > 0) { | |
| if (isSPMDMode()) { | |
| // We have a combined construct, i.e. `target teams distribute | |
| // parallel for [simd]`. We launch so many blocks so that each thread | |
| // will execute one iteration of the loop; rounded up to the nearest | |
| // integer. However, if that results in too few blocks, we artificially | |
| // reduce the thread count per block to increase the outer parallelism. | |
| auto MinThreads = GenericDevice.getMinThreadsForLowTripCountLoop(); | |
| MinThreads = std::min(MinThreads, EffectiveNumThreads); | |
| // Honor the thread_limit clause; only lower the number of threads. | |
| [[maybe_unused]] auto OldNumThreads = EffectiveNumThreads; | |
| if (LoopTripCount >= DefaultNumBlocks * EffectiveNumThreads || | |
| IsNumThreadsFromUser) { | |
| // Enough parallelism for blocks and threads. | |
| TripCountNumBlocks = ((LoopTripCount - 1) / EffectiveNumThreads) + 1; | |
| assert(IsNumThreadsFromUser || | |
| TripCountNumBlocks >= DefaultNumBlocks && | |
| "Expected sufficient outer parallelism."); | |
| } else if (LoopTripCount >= DefaultNumBlocks * MinThreads) { | |
| // Enough parallelism for blocks, limit threads. | |
| // This case is hard; for now, we force "full warps": | |
| // First, compute a thread count assuming DefaultNumBlocks. | |
| auto NumThreadsDefaultBlocks = | |
| (LoopTripCount + DefaultNumBlocks - 1) / DefaultNumBlocks; | |
| // Now get a power of two that is larger or equal. | |
| auto NumThreadsDefaultBlocksP2 = | |
| llvm::PowerOf2Ceil(NumThreadsDefaultBlocks); | |
| // Do not increase a thread limit given be the user. | |
| EffectiveNumThreads = | |
| std::min(EffectiveNumThreads, uint32_t(NumThreadsDefaultBlocksP2)); | |
| assert(EffectiveNumThreads >= MinThreads && | |
| "Expected sufficient inner parallelism."); | |
| TripCountNumBlocks = ((LoopTripCount - 1) / EffectiveNumThreads) + 1; | |
| } else { | |
| // Not enough parallelism for blocks and threads, limit both. | |
| EffectiveNumThreads = std::min(EffectiveNumThreads, MinThreads); | |
| TripCountNumBlocks = ((LoopTripCount - 1) / EffectiveNumThreads) + 1; | |
| } | |
| assert(EffectiveNumThreads * TripCountNumBlocks >= LoopTripCount && | |
| "Expected sufficient parallelism"); | |
| assert(OldNumThreads >= EffectiveNumThreads && | |
| "Number of threads cannot be increased!"); | |
| } else { | |
| assert((isGenericMode() || isGenericSPMDMode()) && | |
| "Unexpected execution mode!"); | |
| // If we reach this point, then we have a non-combined construct, i.e. | |
| // `teams distribute` with a nested `parallel for` and each block is | |
| // assigned one iteration of the `distribute` loop. E.g.: | |
| // | |
| // #pragma omp target teams distribute | |
| // for(...loop_tripcount...) { | |
| // #pragma omp parallel for | |
| // for(...) {} | |
| // } | |
| // | |
| // Threads within a block will execute the iterations of the `parallel` | |
| // loop. | |
| TripCountNumBlocks = LoopTripCount; | |
| } | |
| } | |
| uint32_t PreferredNumBlocks = TripCountNumBlocks; | |
| // If the loops are long running we rather reuse blocks than spawn too many. | |
| if (GenericDevice.getReuseBlocksForHighTripCount()) | |
| PreferredNumBlocks = std::min(TripCountNumBlocks, DefaultNumBlocks); | |
| return std::min(PreferredNumBlocks, | |
| GenericDevice.getBlockLimit(EffectiveNumThreads)); | |
| } | |
| GenericDeviceTy::GenericDeviceTy(GenericPluginTy &Plugin, int32_t DeviceId, | |
| int32_t NumDevices, | |
| const llvm::omp::GV &OMPGridValues) | |
| : Plugin(Plugin), MemoryManager(nullptr), OMP_TeamLimit("OMP_TEAM_LIMIT"), | |
| OMP_NumTeams("OMP_NUM_TEAMS"), | |
| OMP_TeamsThreadLimit("OMP_TEAMS_THREAD_LIMIT"), | |
| OMPX_DebugKind("LIBOMPTARGET_DEVICE_RTL_DEBUG"), | |
| // Do not initialize the following two envars since they depend on the | |
| // device initialization. These cannot be consulted until the device is | |
| // initialized correctly. We initialize them in GenericDeviceTy::init(). | |
| OMPX_TargetStackSize(), OMPX_TargetHeapSize(), | |
| // By default, the initial number of streams and events is 1. | |
| OMPX_InitialNumStreams("LIBOMPTARGET_NUM_INITIAL_STREAMS", 1), | |
| OMPX_InitialNumEvents("LIBOMPTARGET_NUM_INITIAL_EVENTS", 1), | |
| DeviceId(DeviceId), GridValues(OMPGridValues), | |
| PeerAccesses(NumDevices, PeerAccessState::PENDING), PeerAccessesLock(), | |
| PinnedAllocs(*this), RPCServer(nullptr) { | |
| // Conservative fall-back to the plugin's device uid for the case that no real | |
| // vendor (u)uid will become available later. | |
| setDeviceUidFromVendorUid(std::to_string(static_cast<uint64_t>(DeviceId))); | |
| #ifdef OMPT_SUPPORT | |
| OmptInitialized.store(false); | |
| // Bind the callbacks to this device's member functions | |
| #define bindOmptCallback(Name, Type, Code) \ | |
| if (ompt::Initialized && ompt::lookupCallbackByCode) { \ | |
| ompt::lookupCallbackByCode((ompt_callbacks_t)(Code), \ | |
| ((ompt_callback_t *)&(Name##_fn))); \ | |
| ODBG(OLDT_Tool) << "OMPT: class bound " << #Name << "=" \ | |
| << ((void *)(uint64_t)Name##_fn); \ | |
| } | |
| FOREACH_OMPT_DEVICE_EVENT(bindOmptCallback); | |
| #undef bindOmptCallback | |
| #endif | |
| // Envar that indicates whether mapped host buffers should be locked | |
| // automatically. The possible values are boolean (on/off) and a special: | |
| // off: Mapped host buffers are not locked. | |
| // on: Mapped host buffers are locked in a best-effort approach. | |
| // Failure to lock the buffers are silent. | |
| // mandatory: Mapped host buffers are always locked and failures to lock | |
| // a buffer results in a fatal error. | |
| StringEnvar OMPX_LockMappedBuffers("LIBOMPTARGET_LOCK_MAPPED_HOST_BUFFERS", | |
| "off"); | |
| bool Enabled; | |
| if (StringParser::parse(OMPX_LockMappedBuffers.get().data(), Enabled)) { | |
| // Parsed as a boolean value. Enable the feature if necessary. | |
| LockMappedBuffers = Enabled; | |
| IgnoreLockMappedFailures = true; | |
| } else if (OMPX_LockMappedBuffers.get() == "mandatory") { | |
| // Enable the feature and failures are fatal. | |
| LockMappedBuffers = true; | |
| IgnoreLockMappedFailures = false; | |
| } else { | |
| // Disable by default. | |
| ODBG(OLDT_Alloc) << "Invalid value LIBOMPTARGET_LOCK_MAPPED_HOST_BUFFERS=" | |
| << OMPX_LockMappedBuffers.get(); | |
| LockMappedBuffers = false; | |
| } | |
| } | |
| Error GenericDeviceTy::init(GenericPluginTy &Plugin) { | |
| if (auto Err = initImpl(Plugin)) | |
| return Err; | |
| #ifdef OMPT_SUPPORT | |
| if (ompt::Initialized) { | |
| bool ExpectedStatus = false; | |
| if (OmptInitialized.compare_exchange_strong(ExpectedStatus, true)) | |
| performOmptCallback(device_initialize, Plugin.getUserId(DeviceId), | |
| /*type=*/getComputeUnitKind().c_str(), | |
| /*device=*/reinterpret_cast<ompt_device_t *>(this), | |
| /*lookup=*/ompt::lookupCallbackByName, | |
| /*documentation=*/nullptr); | |
| } | |
| #endif | |
| // Read and reinitialize the envars that depend on the device initialization. | |
| // Notice these two envars may change the stack size and heap size of the | |
| // device, so they need the device properly initialized. | |
| auto StackSizeEnvarOrErr = UInt64Envar::create( | |
| "LIBOMPTARGET_STACK_SIZE", | |
| [this](uint64_t &V) -> Error { return getDeviceStackSize(V); }, | |
| [this](uint64_t V) -> Error { return setDeviceStackSize(V); }); | |
| if (!StackSizeEnvarOrErr) | |
| return StackSizeEnvarOrErr.takeError(); | |
| OMPX_TargetStackSize = std::move(*StackSizeEnvarOrErr); | |
| if (hasDeviceHeapSize()) { | |
| auto HeapSizeEnvarOrErr = UInt64Envar::create( | |
| "LIBOMPTARGET_HEAP_SIZE", | |
| [this](uint64_t &V) -> Error { return getDeviceHeapSize(V); }, | |
| [this](uint64_t V) -> Error { return setDeviceHeapSize(V); }); | |
| if (!HeapSizeEnvarOrErr) | |
| return HeapSizeEnvarOrErr.takeError(); | |
| OMPX_TargetHeapSize = std::move(*HeapSizeEnvarOrErr); | |
| } | |
| // Update the maximum number of teams and threads after the device | |
| // initialization sets the corresponding hardware limit. | |
| if (OMP_NumTeams > 0) | |
| GridValues.GV_Max_Teams = | |
| std::min(GridValues.GV_Max_Teams, uint32_t(OMP_NumTeams)); | |
| if (OMP_TeamsThreadLimit > 0) | |
| GridValues.GV_Max_WG_Size = | |
| std::min(GridValues.GV_Max_WG_Size, uint32_t(OMP_TeamsThreadLimit)); | |
| // Enable the memory manager if required. | |
| auto [ThresholdMM, EnableMM] = MemoryManagerTy::getSizeThresholdFromEnv(); | |
| if (EnableMM) { | |
| if (ThresholdMM == 0) | |
| ThresholdMM = getMemoryManagerSizeThreshold(); | |
| MemoryManager = new MemoryManagerTy(*this, ThresholdMM); | |
| } | |
| return Plugin::success(); | |
| } | |
| Error GenericDeviceTy::unloadBinary(DeviceImageTy *Image) { | |
| GenericGlobalHandlerTy &Handler = Plugin.getGlobalHandler(); | |
| auto ProfOrErr = Handler.readProfilingGlobals(*this, *Image); | |
| if (!ProfOrErr) | |
| return ProfOrErr.takeError(); | |
| if (!ProfOrErr->empty()) { | |
| // Dump out profdata | |
| if ((OMPX_DebugKind.get() & uint32_t(DeviceDebugKind::PGODump)) == | |
| uint32_t(DeviceDebugKind::PGODump)) | |
| ProfOrErr->dump(); | |
| // Write data to profiling file | |
| if (auto Err = ProfOrErr->write()) | |
| return Err; | |
| } | |
| return unloadBinaryImpl(Image); | |
| } | |
| Error GenericDeviceTy::deinit(GenericPluginTy &Plugin) { | |
| // Run the global destructors first in case they required the RPC server. | |
| for (auto &I : LoadedImages) { | |
| if (auto Err = callGlobalDestructors(Plugin, *I)) | |
| return Err; | |
| } | |
| if (RPCServer) { | |
| if (auto Err = RPCServer->deinitDevice(*this)) | |
| return Err; | |
| RPCServer = nullptr; | |
| } | |
| for (auto &I : LoadedImages) | |
| if (auto Err = unloadBinary(I)) | |
| return Err; | |
| LoadedImages.clear(); | |
| // Delete the memory manager before deinitializing the device. Otherwise, | |
| // we may delete device allocations after the device is deinitialized. | |
| if (MemoryManager) | |
| delete MemoryManager; | |
| MemoryManager = nullptr; | |
| if (RecordReplay) { | |
| if (auto Err = RecordReplay->deinit()) | |
| return Err; | |
| delete RecordReplay; | |
| RecordReplay = nullptr; | |
| } | |
| #ifdef OMPT_SUPPORT | |
| if (ompt::Initialized) { | |
| bool ExpectedStatus = true; | |
| if (OmptInitialized.compare_exchange_strong(ExpectedStatus, false)) | |
| performOmptCallback(device_finalize, Plugin.getUserId(DeviceId)); | |
| } | |
| #endif | |
| return deinitImpl(); | |
| } | |
| Expected<DeviceImageTy *> GenericDeviceTy::loadBinary(GenericPluginTy &Plugin, | |
| StringRef InputTgtImage) { | |
| ODBG(OLDT_Init) << "Load data from image " | |
| << static_cast<const void *>(InputTgtImage.bytes_begin()); | |
| std::unique_ptr<MemoryBuffer> Buffer; | |
| if (identify_magic(InputTgtImage) == file_magic::bitcode) { | |
| auto CompiledImageOrErr = Plugin.getJIT().process(InputTgtImage, *this); | |
| if (!CompiledImageOrErr) { | |
| return Plugin::error(ErrorCode::COMPILE_FAILURE, | |
| CompiledImageOrErr.takeError(), | |
| "failure to jit IR image"); | |
| } | |
| Buffer = std::move(*CompiledImageOrErr); | |
| } else { | |
| Buffer = MemoryBuffer::getMemBufferCopy(InputTgtImage); | |
| } | |
| // Load the binary and allocate the image object. Use the next available id | |
| // for the image id, which is the number of previously loaded images. | |
| auto ImageOrErr = loadBinaryImpl(std::move(Buffer), LoadedImages.size()); | |
| if (!ImageOrErr) | |
| return ImageOrErr.takeError(); | |
| DeviceImageTy *Image = *ImageOrErr; | |
| if (identify_magic(InputTgtImage) == file_magic::bitcode) | |
| Image->setIRImage(MemoryBuffer::getMemBufferCopy(InputTgtImage)); | |
| // Add the image to list. | |
| LoadedImages.push_back(Image); | |
| if (auto Err = setupRPCServer(Plugin, *Image)) | |
| return std::move(Err); | |
| #ifdef OMPT_SUPPORT | |
| if (ompt::Initialized) { | |
| size_t Bytes = InputTgtImage.size(); | |
| performOmptCallback( | |
| device_load, Plugin.getUserId(DeviceId), | |
| /*FileName=*/nullptr, /*FileOffset=*/0, /*VmaInFile=*/nullptr, | |
| /*ImgSize=*/Bytes, | |
| /*HostAddr=*/const_cast<unsigned char *>(InputTgtImage.bytes_begin()), | |
| /*DeviceAddr=*/nullptr, /* FIXME: ModuleId */ 0); | |
| } | |
| #endif | |
| // Call any global constructors present on the device. | |
| if (auto Err = callGlobalConstructors(Plugin, *Image)) | |
| return std::move(Err); | |
| // Return the pointer to the table of entries. | |
| return Image; | |
| } | |
| Error GenericDeviceTy::setupRPCServer(GenericPluginTy &Plugin, | |
| DeviceImageTy &Image) { | |
| // The plugin either does not need an RPC server or it is unavailable. | |
| if (!shouldSetupRPCServer()) | |
| return Plugin::success(); | |
| // Check if this device needs to run an RPC server. | |
| RPCServerTy &Server = Plugin.getRPCServer(); | |
| auto UsingOrErr = | |
| Server.isDeviceUsingRPC(*this, Plugin.getGlobalHandler(), Image); | |
| if (!UsingOrErr) | |
| return UsingOrErr.takeError(); | |
| if (!UsingOrErr.get()) | |
| return Plugin::success(); | |
| if (auto Err = Server.initDevice(*this, Plugin.getGlobalHandler(), Image)) | |
| return Err; | |
| if (auto Err = Server.startThread()) | |
| return Err; | |
| RPCServer = &Server; | |
| ODBG(OLDT_Init) << "Running an RPC server on device " << getDeviceId(); | |
| return Plugin::success(); | |
| } | |
| Error PinnedAllocationMapTy::insertEntry(void *HstPtr, void *DevAccessiblePtr, | |
| size_t Size, bool ExternallyLocked) { | |
| // Insert the new entry into the map. | |
| auto Res = Allocs.insert({HstPtr, DevAccessiblePtr, Size, ExternallyLocked}); | |
| if (!Res.second) | |
| return Plugin::error(ErrorCode::INVALID_ARGUMENT, | |
| "cannot insert locked buffer entry"); | |
| // Check whether the next entry overlaps with the inserted entry. | |
| auto It = std::next(Res.first); | |
| if (It == Allocs.end()) | |
| return Plugin::success(); | |
| const EntryTy *NextEntry = &(*It); | |
| if (intersects(NextEntry->HstPtr, NextEntry->Size, HstPtr, Size)) | |
| return Plugin::error(ErrorCode::INVALID_ARGUMENT, | |
| "partial overlapping not allowed in locked buffers"); | |
| return Plugin::success(); | |
| } | |
| Error PinnedAllocationMapTy::eraseEntry(const EntryTy &Entry) { | |
| // Erase the existing entry. Notice this requires an additional map lookup, | |
| // but this should not be a performance issue. Using iterators would make | |
| // the code more difficult to read. | |
| size_t Erased = Allocs.erase({Entry.HstPtr}); | |
| if (!Erased) | |
| return Plugin::error(ErrorCode::INVALID_ARGUMENT, | |
| "cannot erase locked buffer entry"); | |
| return Plugin::success(); | |
| } | |
| Error PinnedAllocationMapTy::registerEntryUse(const EntryTy &Entry, | |
| void *HstPtr, size_t Size) { | |
| if (!contains(Entry.HstPtr, Entry.Size, HstPtr, Size)) | |
| return Plugin::error(ErrorCode::INVALID_ARGUMENT, | |
| "partial overlapping not allowed in locked buffers"); | |
| ++Entry.References; | |
| return Plugin::success(); | |
| } | |
| Expected<bool> PinnedAllocationMapTy::unregisterEntryUse(const EntryTy &Entry) { | |
| if (Entry.References == 0) | |
| return Plugin::error(ErrorCode::INVALID_ARGUMENT, | |
| "invalid number of references"); | |
| // Return whether this was the last user. | |
| return (--Entry.References == 0); | |
| } | |
| Error PinnedAllocationMapTy::registerHostBuffer(void *HstPtr, | |
| void *DevAccessiblePtr, | |
| size_t Size) { | |
| assert(HstPtr && "Invalid pointer"); | |
| assert(DevAccessiblePtr && "Invalid pointer"); | |
| assert(Size && "Invalid size"); | |
| std::lock_guard<std::shared_mutex> Lock(Mutex); | |
| // No pinned allocation should intersect. | |
| const EntryTy *Entry = findIntersecting(HstPtr); | |
| if (Entry) | |
| return Plugin::error(ErrorCode::INVALID_ARGUMENT, | |
| "cannot insert entry due to an existing one"); | |
| // Now insert the new entry. | |
| return insertEntry(HstPtr, DevAccessiblePtr, Size); | |
| } | |
| Error PinnedAllocationMapTy::unregisterHostBuffer(void *HstPtr) { | |
| assert(HstPtr && "Invalid pointer"); | |
| std::lock_guard<std::shared_mutex> Lock(Mutex); | |
| const EntryTy *Entry = findIntersecting(HstPtr); | |
| if (!Entry) | |
| return Plugin::error(ErrorCode::INVALID_ARGUMENT, | |
| "cannot find locked buffer"); | |
| // The address in the entry should be the same we are unregistering. | |
| if (Entry->HstPtr != HstPtr) | |
| return Plugin::error(ErrorCode::INVALID_ARGUMENT, | |
| "unexpected host pointer in locked buffer entry"); | |
| // Unregister from the entry. | |
| auto LastUseOrErr = unregisterEntryUse(*Entry); | |
| if (!LastUseOrErr) | |
| return LastUseOrErr.takeError(); | |
| // There should be no other references to the pinned allocation. | |
| if (!(*LastUseOrErr)) | |
| return Plugin::error(ErrorCode::INVALID_ARGUMENT, | |
| "the locked buffer is still being used"); | |
| // Erase the entry from the map. | |
| return eraseEntry(*Entry); | |
| } | |
| Expected<void *> PinnedAllocationMapTy::registerMemory(void *HstPtr, | |
| size_t Size, | |
| bool LockMemory) { | |
| assert(HstPtr && "Invalid pointer"); | |
| assert(Size && "Invalid size"); | |
| std::lock_guard<std::shared_mutex> Lock(Mutex); | |
| const EntryTy *Entry = findIntersecting(HstPtr); | |
| if (Entry) { | |
| // An already registered intersecting buffer was found. Register a new use. | |
| if (auto Err = registerEntryUse(*Entry, HstPtr, Size)) | |
| return std::move(Err); | |
| // Return the device accessible pointer with the correct offset. | |
| return utils::advancePtr(Entry->DevAccessiblePtr, | |
| utils::getPtrDiff(HstPtr, Entry->HstPtr)); | |
| } | |
| size_t BaseSize; | |
| void *BaseHstPtr, *BaseDevAccessiblePtr; | |
| // Check if it was externally pinned by a vendor-specific API. | |
| auto IsPinnedOrErr = Device.isPinnedPtrImpl(HstPtr, BaseHstPtr, | |
| BaseDevAccessiblePtr, BaseSize); | |
| if (!IsPinnedOrErr) | |
| return IsPinnedOrErr.takeError(); | |
| // If pinned, just insert the entry representing the whole pinned buffer. | |
| if (*IsPinnedOrErr) { | |
| if (auto Err = insertEntry(BaseHstPtr, BaseDevAccessiblePtr, BaseSize, | |
| /*Externallylocked=*/true)) | |
| return std::move(Err); | |
| return BaseDevAccessiblePtr; | |
| } | |
| // Not externally pinned. Do nothing if locking of mapped buffers is disabled. | |
| if (!LockMemory) | |
| return nullptr; | |
| // No intersecting registered allocation found in the map. First, lock the | |
| // host buffer and retrieve the device accessible pointer. | |
| auto DevAccessiblePtrOrErr = Device.dataLockImpl(HstPtr, Size); | |
| if (!DevAccessiblePtrOrErr) | |
| return DevAccessiblePtrOrErr.takeError(); | |
| // Now insert the new entry into the map. | |
| if (auto Err = insertEntry(HstPtr, *DevAccessiblePtrOrErr, Size)) | |
| return std::move(Err); | |
| // Return the device accessible pointer. | |
| return *DevAccessiblePtrOrErr; | |
| } | |
| Error PinnedAllocationMapTy::unregisterMemory(void *HstPtr, bool UnlockMemory) { | |
| assert(HstPtr && "Invalid pointer"); | |
| std::lock_guard<std::shared_mutex> Lock(Mutex); | |
| const EntryTy *Entry = findIntersecting(HstPtr); | |
| // No entry but automatic locking of mapped buffers is disabled, so | |
| // nothing to do. | |
| if (!Entry && !UnlockMemory) | |
| return Plugin::success(); | |
| if (!Entry) | |
| return Plugin::error(ErrorCode::INVALID_ARGUMENT, | |
| "cannot find locked buffer"); | |
| // Unregister from the locked buffer. No need to do anything if there are | |
| // others using the allocation. | |
| auto LastUseOrErr = unregisterEntryUse(*Entry); | |
| if (!LastUseOrErr) | |
| return LastUseOrErr.takeError(); | |
| // No need to do anything if there are others using the allocation. | |
| if (!(*LastUseOrErr)) | |
| return Plugin::success(); | |
| // This was the last user of the allocation. Unlock the original locked buffer | |
| // if it was locked by the plugin. Do not unlock it if it was locked by an | |
| // external entity. Unlock the buffer using the host pointer of the entry. | |
| if (!Entry->ExternallyLocked) | |
| if (auto Err = Device.dataUnlockImpl(Entry->HstPtr)) | |
| return Err; | |
| // Erase the entry from the map. | |
| return eraseEntry(*Entry); | |
| } | |
| Error GenericDeviceTy::synchronize(__tgt_async_info *AsyncInfo, | |
| bool ReleaseQueue) { | |
| if (!AsyncInfo) | |
| return Plugin::error(ErrorCode::INVALID_ARGUMENT, | |
| "invalid async info queue"); | |
| SmallVector<void *> AllocsToDelete{}; | |
| { | |
| std::lock_guard<std::mutex> AllocationGuard{AsyncInfo->Mutex}; | |
| // This can be false when no work has been added to the AsyncInfo. In which | |
| // case, the device has nothing to synchronize. | |
| if (AsyncInfo->Queue) | |
| if (auto Err = synchronizeImpl(*AsyncInfo, ReleaseQueue)) | |
| return Err; | |
| std::swap(AllocsToDelete, AsyncInfo->AssociatedAllocations); | |
| } | |
| for (auto *Ptr : AllocsToDelete) | |
| if (auto Err = dataDelete(Ptr, TargetAllocTy::TARGET_ALLOC_DEVICE)) | |
| return Err; | |
| return Plugin::success(); | |
| } | |
| Error GenericDeviceTy::queryAsync(__tgt_async_info *AsyncInfo, | |
| bool ReleaseQueue, | |
| bool *IsQueueWorkCompleted) { | |
| if (!AsyncInfo || !AsyncInfo->Queue) | |
| return Plugin::error(ErrorCode::INVALID_ARGUMENT, | |
| "invalid async info queue"); | |
| return queryAsyncImpl(*AsyncInfo, ReleaseQueue, IsQueueWorkCompleted); | |
| } | |
| Error GenericDeviceTy::memoryVAMap(void **Addr, void *VAddr, size_t *RSize) { | |
| return Plugin::error(ErrorCode::UNSUPPORTED, | |
| "device does not support VA Management"); | |
| } | |
| Error GenericDeviceTy::memoryVAUnMap(void *VAddr, size_t Size) { | |
| return Plugin::error(ErrorCode::UNSUPPORTED, | |
| "device does not support VA Management"); | |
| } | |
| Error GenericDeviceTy::getDeviceMemorySize(uint64_t &DSize) { | |
| return Plugin::error( | |
| ErrorCode::UNIMPLEMENTED, | |
| "missing getDeviceMemorySize implementation (required by RR-heuristic"); | |
| } | |
| Expected<void *> GenericDeviceTy::dataAlloc(int64_t Size, void *HostPtr, | |
| TargetAllocTy Kind, | |
| size_t Alignment) { | |
| void *Alloc = nullptr; | |
| // TODO Check alignment. | |
| if (RecordReplay && RecordReplay->isRecordingOrReplaying()) | |
| return RecordReplay->allocate(Size); | |
| switch (Kind) { | |
| case TARGET_ALLOC_DEFAULT: | |
| case TARGET_ALLOC_DEVICE: | |
| if (MemoryManager) { | |
| auto AllocOrErr = MemoryManager->allocate(Size, HostPtr, Alignment); | |
| if (!AllocOrErr) | |
| return AllocOrErr.takeError(); | |
| Alloc = *AllocOrErr; | |
| if (!Alloc) | |
| return Plugin::error(ErrorCode::OUT_OF_RESOURCES, | |
| "failed to allocate from memory manager"); | |
| break; | |
| } | |
| [[fallthrough]]; | |
| case TARGET_ALLOC_HOST: | |
| case TARGET_ALLOC_SHARED: { | |
| auto AllocOrErr = allocate(Size, HostPtr, Kind, Alignment); | |
| if (!AllocOrErr) | |
| return AllocOrErr.takeError(); | |
| Alloc = *AllocOrErr; | |
| if (!Alloc) | |
| return Plugin::error(ErrorCode::OUT_OF_RESOURCES, | |
| "failed to allocate from device allocator"); | |
| } | |
| } | |
| // Report error if the memory manager or the device allocator did not return | |
| // any memory buffer. | |
| if (!Alloc) | |
| return Plugin::error(ErrorCode::UNIMPLEMENTED, | |
| "invalid target data allocation kind or requested " | |
| "allocator not implemented yet"); | |
| // Register allocated buffer as pinned memory if the type is host memory. | |
| if (Kind == TARGET_ALLOC_HOST) | |
| if (auto Err = PinnedAllocs.registerHostBuffer(Alloc, Alloc, Size)) | |
| return std::move(Err); | |
| // Keep track of the allocation stack if we track allocation traces. | |
| if (OMPX_TrackAllocationTraces) { | |
| std::string StackTrace; | |
| llvm::raw_string_ostream OS(StackTrace); | |
| llvm::sys::PrintStackTrace(OS); | |
| AllocationTraceInfoTy *ATI = new AllocationTraceInfoTy(); | |
| ATI->AllocationTrace = std::move(StackTrace); | |
| ATI->DevicePtr = Alloc; | |
| ATI->HostPtr = HostPtr; | |
| ATI->Size = Size; | |
| ATI->Kind = Kind; | |
| auto AllocationTraceMap = AllocationTraces.getExclusiveAccessor(); | |
| auto *&MapATI = (*AllocationTraceMap)[Alloc]; | |
| ATI->LastAllocationInfo = MapATI; | |
| MapATI = ATI; | |
| } | |
| return Alloc; | |
| } | |
| Error GenericDeviceTy::dataDelete(void *TgtPtr, TargetAllocTy Kind) { | |
| // Free is a noop when recording or replaying. | |
| if (RecordReplay && RecordReplay->isRecordingOrReplaying()) | |
| return RecordReplay->deallocate(TgtPtr); | |
| // Keep track of the deallocation stack if we track allocation traces. | |
| if (OMPX_TrackAllocationTraces) { | |
| AllocationTraceInfoTy *ATI = nullptr; | |
| { | |
| auto AllocationTraceMap = AllocationTraces.getExclusiveAccessor(); | |
| ATI = (*AllocationTraceMap)[TgtPtr]; | |
| } | |
| std::string StackTrace; | |
| llvm::raw_string_ostream OS(StackTrace); | |
| llvm::sys::PrintStackTrace(OS); | |
| if (!ATI) | |
| ErrorReporter::reportDeallocationOfNonAllocatedPtr(TgtPtr, Kind, ATI, | |
| StackTrace); | |
| // ATI is not null, thus we can lock it to inspect and modify it further. | |
| std::lock_guard<std::mutex> LG(ATI->Lock); | |
| if (!ATI->DeallocationTrace.empty()) | |
| ErrorReporter::reportDeallocationOfDeallocatedPtr(TgtPtr, Kind, ATI, | |
| StackTrace); | |
| if (ATI->Kind != Kind) | |
| ErrorReporter::reportDeallocationOfWrongPtrKind(TgtPtr, Kind, ATI, | |
| StackTrace); | |
| ATI->DeallocationTrace = StackTrace; | |
| } | |
| switch (Kind) { | |
| case TARGET_ALLOC_DEFAULT: | |
| case TARGET_ALLOC_DEVICE: | |
| if (MemoryManager) { | |
| if (auto Err = MemoryManager->free(TgtPtr)) | |
| return Err; | |
| break; | |
| } | |
| [[fallthrough]]; | |
| case TARGET_ALLOC_HOST: | |
| case TARGET_ALLOC_SHARED: | |
| if (auto Err = free(TgtPtr, Kind)) | |
| return Err; | |
| } | |
| // Unregister deallocated pinned memory buffer if the type is host memory. | |
| if (Kind == TARGET_ALLOC_HOST) | |
| if (auto Err = PinnedAllocs.unregisterHostBuffer(TgtPtr)) | |
| return Err; | |
| return Plugin::success(); | |
| } | |
| Error GenericDeviceTy::dataSubmit(void *TgtPtr, const void *HstPtr, | |
| int64_t Size, __tgt_async_info *AsyncInfo) { | |
| AsyncInfoWrapperTy AsyncInfoWrapper(*this, AsyncInfo); | |
| auto Err = dataSubmitImpl(TgtPtr, HstPtr, Size, AsyncInfoWrapper); | |
| AsyncInfoWrapper.finalize(Err); | |
| return Err; | |
| } | |
| Error GenericDeviceTy::dataRetrieve(void *HstPtr, const void *TgtPtr, | |
| int64_t Size, __tgt_async_info *AsyncInfo) { | |
| AsyncInfoWrapperTy AsyncInfoWrapper(*this, AsyncInfo); | |
| auto Err = dataRetrieveImpl(HstPtr, TgtPtr, Size, AsyncInfoWrapper); | |
| AsyncInfoWrapper.finalize(Err); | |
| return Err; | |
| } | |
| Error GenericDeviceTy::dataExchange(const void *SrcPtr, GenericDeviceTy &DstDev, | |
| void *DstPtr, int64_t Size, | |
| __tgt_async_info *AsyncInfo) { | |
| AsyncInfoWrapperTy AsyncInfoWrapper(*this, AsyncInfo); | |
| auto Err = dataExchangeImpl(SrcPtr, DstDev, DstPtr, Size, AsyncInfoWrapper); | |
| AsyncInfoWrapper.finalize(Err); | |
| return Err; | |
| } | |
| Error GenericDeviceTy::dataFill(void *TgtPtr, const void *PatternPtr, | |
| int64_t PatternSize, int64_t Size, | |
| __tgt_async_info *AsyncInfo) { | |
| AsyncInfoWrapperTy AsyncInfoWrapper(*this, AsyncInfo); | |
| auto Err = | |
| dataFillImpl(TgtPtr, PatternPtr, PatternSize, Size, AsyncInfoWrapper); | |
| AsyncInfoWrapper.finalize(Err); | |
| return Err; | |
| } | |
| Error GenericDeviceTy::dataPrefetch(size_t Count, const void **Mems, | |
| const size_t *Sizes, bool ToHost, | |
| __tgt_async_info *AsyncInfo) { | |
| AsyncInfoWrapperTy AsyncInfoWrapper(*this, AsyncInfo); | |
| auto Err = dataPrefetchImpl(Count, Mems, Sizes, ToHost, AsyncInfoWrapper); | |
| AsyncInfoWrapper.finalize(Err); | |
| return Err; | |
| } | |
| Error GenericDeviceTy::launchKernel(void *EntryPtr, void **ArgPtrs, | |
| ptrdiff_t *ArgOffsets, | |
| KernelArgsTy &KernelArgs, | |
| KernelExtraArgsTy *KernelExtraArgs, | |
| __tgt_async_info *AsyncInfo) { | |
| AsyncInfoWrapperTy AsyncInfoWrapper(*this, AsyncInfo); | |
| GenericKernelTy &GenericKernel = | |
| *reinterpret_cast<GenericKernelTy *>(EntryPtr); | |
| { | |
| std::string StackTrace; | |
| if (OMPX_TrackNumKernelLaunches) { | |
| llvm::raw_string_ostream OS(StackTrace); | |
| llvm::sys::PrintStackTrace(OS); | |
| } | |
| auto KernelTraceInfoRecord = KernelLaunchTraces.getExclusiveAccessor(); | |
| (*KernelTraceInfoRecord) | |
| .emplace(&GenericKernel, std::move(StackTrace), AsyncInfo); | |
| } | |
| auto Err = GenericKernel.launch(*this, ArgPtrs, ArgOffsets, KernelArgs, | |
| KernelExtraArgs, AsyncInfoWrapper); | |
| AsyncInfoWrapper.finalize(Err); | |
| return Err; | |
| } | |
| Error GenericDeviceTy::initAsyncInfo(__tgt_async_info **AsyncInfoPtr) { | |
| assert(AsyncInfoPtr && "Invalid async info"); | |
| *AsyncInfoPtr = new __tgt_async_info(); | |
| AsyncInfoWrapperTy AsyncInfoWrapper(*this, *AsyncInfoPtr); | |
| auto Err = initAsyncInfoImpl(AsyncInfoWrapper); | |
| AsyncInfoWrapper.finalize(Err); | |
| return Err; | |
| } | |
| Error GenericDeviceTy::enqueueHostCall(void (*Callback)(void *), void *UserData, | |
| __tgt_async_info *AsyncInfo) { | |
| AsyncInfoWrapperTy AsyncInfoWrapper(*this, AsyncInfo); | |
| auto Err = enqueueHostCallImpl(Callback, UserData, AsyncInfoWrapper); | |
| AsyncInfoWrapper.finalize(Err); | |
| return Err; | |
| } | |
| Expected<InfoTreeNode> GenericDeviceTy::obtainInfo() { | |
| auto InfoOrErr = obtainInfoImpl(); | |
| if (InfoOrErr) | |
| InfoOrErr->add("UID", getDeviceUid(), "", DeviceInfo::UID); | |
| return InfoOrErr; | |
| } | |
| Error GenericDeviceTy::printInfo() { | |
| auto InfoOrErr = obtainInfo(); | |
| // Get the vendor-specific info entries describing the device properties. | |
| if (auto Err = InfoOrErr.takeError()) | |
| return Err; | |
| // Print all info entries. | |
| InfoOrErr->print(); | |
| return Plugin::success(); | |
| } | |
| Error GenericDeviceTy::createEvent(void **EventPtrStorage, | |
| bool EnableProfiling) { | |
| return createEventImpl(EventPtrStorage, EnableProfiling); | |
| } | |
| Error GenericDeviceTy::destroyEvent(void *EventPtr, bool EnableProfiling) { | |
| return destroyEventImpl(EventPtr, EnableProfiling); | |
| } | |
| Error GenericDeviceTy::recordEvent(void *EventPtr, __tgt_async_info *AsyncInfo, | |
| bool EnableProfiling) { | |
| AsyncInfoWrapperTy AsyncInfoWrapper(*this, AsyncInfo); | |
| auto Err = recordEventImpl(EventPtr, AsyncInfoWrapper, EnableProfiling); | |
| AsyncInfoWrapper.finalize(Err); | |
| return Err; | |
| } | |
| Error GenericDeviceTy::waitEvent(void *EventPtr, __tgt_async_info *AsyncInfo) { | |
| AsyncInfoWrapperTy AsyncInfoWrapper(*this, AsyncInfo); | |
| auto Err = waitEventImpl(EventPtr, AsyncInfoWrapper); | |
| AsyncInfoWrapper.finalize(Err); | |
| return Err; | |
| } | |
| Expected<bool> GenericDeviceTy::hasPendingWork(__tgt_async_info *AsyncInfo) { | |
| AsyncInfoWrapperTy AsyncInfoWrapper(*this, AsyncInfo); | |
| auto Res = hasPendingWorkImpl(AsyncInfoWrapper); | |
| if (auto Err = Res.takeError()) { | |
| AsyncInfoWrapper.finalize(Err); | |
| return Err; | |
| } | |
| auto Err = Plugin::success(); | |
| AsyncInfoWrapper.finalize(Err); | |
| if (Err) | |
| return Err; | |
| return Res; | |
| } | |
| Expected<bool> GenericDeviceTy::isEventComplete(void *Event, | |
| __tgt_async_info *AsyncInfo) { | |
| AsyncInfoWrapperTy AsyncInfoWrapper(*this, AsyncInfo); | |
| auto Res = isEventCompleteImpl(Event, AsyncInfoWrapper); | |
| if (auto Err = Res.takeError()) { | |
| AsyncInfoWrapper.finalize(Err); | |
| return Err; | |
| } | |
| auto Err = Plugin::success(); | |
| AsyncInfoWrapper.finalize(Err); | |
| if (Err) | |
| return Err; | |
| return Res; | |
| } | |
| Error GenericDeviceTy::syncEvent(void *EventPtr) { | |
| return syncEventImpl(EventPtr); | |
| } | |
| Expected<float> GenericDeviceTy::getEventElapsedTime(void *StartEventPtr, | |
| void *EndEventPtr) { | |
| return getEventElapsedTimeImpl(StartEventPtr, EndEventPtr); | |
| } | |
| bool GenericDeviceTy::useAutoZeroCopy() { return useAutoZeroCopyImpl(); } | |
| Expected<bool> GenericDeviceTy::isAccessiblePtr(const void *Ptr, size_t Size) { | |
| return isAccessiblePtrImpl(Ptr, Size); | |
| } | |
| void GenericDeviceTy::setDeviceUidFromVendorUid(StringRef VendorUid) { | |
| DeviceUid = std::string(Plugin.getName()) + "-" + std::string(VendorUid); | |
| } | |
| Error GenericPluginTy::init() { | |
| if (Initialized) | |
| return Plugin::success(); | |
| auto NumDevicesOrErr = initImpl(); | |
| if (!NumDevicesOrErr) | |
| return NumDevicesOrErr.takeError(); | |
| Initialized = true; | |
| NumDevices = *NumDevicesOrErr; | |
| if (NumDevices == 0) | |
| return Plugin::success(); | |
| assert(Devices.size() == 0 && "Plugin already initialized"); | |
| Devices.resize(NumDevices, nullptr); | |
| GlobalHandler = createGlobalHandler(); | |
| assert(GlobalHandler && "Invalid global handler"); | |
| RPCServer = new RPCServerTy(*this); | |
| assert(RPCServer && "Invalid RPC server"); | |
| return Plugin::success(); | |
| } | |
| Error GenericPluginTy::deinit() { | |
| assert(Initialized && "Plugin was not initialized!"); | |
| // Deinitialize all active devices. | |
| for (int32_t DeviceId = 0; DeviceId < NumDevices; ++DeviceId) { | |
| if (Devices[DeviceId]) { | |
| if (auto Err = deinitDevice(DeviceId)) | |
| return Err; | |
| } | |
| assert(!Devices[DeviceId] && "Device was not deinitialized"); | |
| } | |
| // There is no global handler if no device is available. | |
| if (GlobalHandler) | |
| delete GlobalHandler; | |
| if (RPCServer) { | |
| if (Error Err = RPCServer->shutDown(*this)) | |
| return Err; | |
| delete RPCServer; | |
| } | |
| // Perform last deinitializations on the plugin. | |
| if (Error Err = deinitImpl()) | |
| return Err; | |
| Initialized = false; | |
| return Plugin::success(); | |
| } | |
| Error GenericPluginTy::initDevice(int32_t DeviceId) { | |
| assert(!Devices[DeviceId] && "Device already initialized"); | |
| // Create the device and save the reference. | |
| GenericDeviceTy *Device = createDevice(*this, DeviceId, NumDevices); | |
| assert(Device && "Invalid device"); | |
| // Save the device reference into the list. | |
| Devices[DeviceId] = Device; | |
| // Initialize the device and its resources. | |
| return Device->init(*this); | |
| } | |
| Error GenericPluginTy::deinitDevice(int32_t DeviceId) { | |
| // The device may be already deinitialized. | |
| if (Devices[DeviceId] == nullptr) | |
| return Plugin::success(); | |
| // Deinitialize the device and release its resources. | |
| if (auto Err = Devices[DeviceId]->deinit(*this)) | |
| return Err; | |
| // Delete the device and invalidate its reference. | |
| delete Devices[DeviceId]; | |
| Devices[DeviceId] = nullptr; | |
| return Plugin::success(); | |
| } | |
| Expected<bool> GenericPluginTy::checkELFImage(StringRef Image) const { | |
| // First check if this image is a regular ELF file. | |
| if (!utils::elf::isELF(Image)) | |
| return false; | |
| // Check if this image is an ELF with a matching machine value. | |
| auto MachineOrErr = utils::elf::checkMachine(Image, getMagicElfBits()); | |
| if (!MachineOrErr) | |
| return MachineOrErr.takeError(); | |
| return MachineOrErr; | |
| } | |
| Expected<bool> GenericPluginTy::checkBitcodeImage(StringRef Image) const { | |
| if (identify_magic(Image) != file_magic::bitcode) | |
| return false; | |
| LLVMContext Context; | |
| auto ModuleOrErr = getLazyBitcodeModule(MemoryBufferRef(Image, ""), Context, | |
| /*ShouldLazyLoadMetadata=*/true); | |
| if (!ModuleOrErr) | |
| return ModuleOrErr.takeError(); | |
| Module &M = **ModuleOrErr; | |
| return M.getTargetTriple().getArch() == getTripleArch(); | |
| } | |
| int32_t GenericPluginTy::is_initialized() const { return Initialized; } | |
| int32_t GenericPluginTy::isPluginCompatible(StringRef Image) { | |
| auto HandleError = [&](Error Err) -> bool { | |
| std::string ErrStr = toString(std::move(Err)); | |
| ODBG(OLDT_Init) << "Failure to check validity of image " | |
| << static_cast<const void *>(Image.data()) << ": " | |
| << ErrStr; | |
| return false; | |
| }; | |
| switch (identify_magic(Image)) { | |
| case file_magic::elf: | |
| case file_magic::elf_relocatable: | |
| case file_magic::elf_executable: | |
| case file_magic::elf_shared_object: | |
| case file_magic::elf_core: { | |
| auto MatchOrErr = checkELFImage(Image); | |
| if (Error Err = MatchOrErr.takeError()) | |
| return HandleError(std::move(Err)); | |
| return *MatchOrErr; | |
| } | |
| case file_magic::bitcode: { | |
| auto MatchOrErr = checkBitcodeImage(Image); | |
| if (Error Err = MatchOrErr.takeError()) | |
| return HandleError(std::move(Err)); | |
| return *MatchOrErr; | |
| } | |
| default: | |
| auto MatchOrErr = isImageCompatible(Image); | |
| if (Error Err = MatchOrErr.takeError()) | |
| return HandleError(std::move(Err)); | |
| return *MatchOrErr; | |
| } | |
| } | |
| int32_t GenericPluginTy::isDeviceCompatible(int32_t DeviceId, StringRef Image) { | |
| auto HandleError = [&](Error Err) -> bool { | |
| std::string ErrStr = toString(std::move(Err)); | |
| ODBG(OLDT_Init) << "Failure to check validity of image " | |
| << static_cast<const void *>(Image.data()) << ": " | |
| << ErrStr; | |
| return false; | |
| }; | |
| switch (identify_magic(Image)) { | |
| case file_magic::elf: | |
| case file_magic::elf_relocatable: | |
| case file_magic::elf_executable: | |
| case file_magic::elf_shared_object: | |
| case file_magic::elf_core: { | |
| auto MatchOrErr = checkELFImage(Image); | |
| if (Error Err = MatchOrErr.takeError()) | |
| return HandleError(std::move(Err)); | |
| if (!*MatchOrErr) | |
| return false; | |
| // Perform plugin-dependent checks for the specific architecture if needed. | |
| auto CompatibleOrErr = isELFCompatible(DeviceId, Image); | |
| if (Error Err = CompatibleOrErr.takeError()) | |
| return HandleError(std::move(Err)); | |
| return *CompatibleOrErr; | |
| } | |
| case file_magic::bitcode: { | |
| auto MatchOrErr = checkBitcodeImage(Image); | |
| if (Error Err = MatchOrErr.takeError()) | |
| return HandleError(std::move(Err)); | |
| return *MatchOrErr; | |
| } | |
| default: | |
| auto MatchOrErr = isImageCompatible(DeviceId, Image); | |
| if (Error Err = MatchOrErr.takeError()) | |
| return HandleError(std::move(Err)); | |
| return *MatchOrErr; | |
| } | |
| } | |
| int32_t GenericPluginTy::is_device_initialized(int32_t DeviceId) const { | |
| return isValidDeviceId(DeviceId) && Devices[DeviceId] != nullptr; | |
| } | |
| int32_t GenericPluginTy::init_device(int32_t DeviceId) { | |
| auto Err = initDevice(DeviceId); | |
| if (Err) { | |
| REPORT() << "Failure to initialize device " << DeviceId << ": " | |
| << toString(std::move(Err)); | |
| return OFFLOAD_FAIL; | |
| } | |
| return OFFLOAD_SUCCESS; | |
| } | |
| int32_t GenericPluginTy::number_of_devices() { return getNumDevices(); } | |
| int32_t GenericPluginTy::is_data_exchangable(int32_t SrcDeviceId, | |
| int32_t DstDeviceId) { | |
| return isDataExchangable(SrcDeviceId, DstDeviceId); | |
| } | |
| int32_t GenericPluginTy::initialize_record_replay( | |
| int32_t DeviceId, int64_t MemorySize, void *VAddr, bool IsRecord, | |
| bool IsNative, bool SaveOutput, bool EmitReport, const char *ReportFilename, | |
| const char *OutputDirPath) { | |
| GenericDeviceTy &Device = getDevice(DeviceId); | |
| if (auto Err = Device.initRecordReplay(MemorySize, VAddr, IsRecord, IsNative, | |
| SaveOutput, EmitReport, ReportFilename, | |
| OutputDirPath)) { | |
| REPORT() << "Failure to initialize RR with " << MemorySize | |
| << " bytes on device " << DeviceId << ": " | |
| << toString(std::move(Err)); | |
| return OFFLOAD_FAIL; | |
| } | |
| return OFFLOAD_SUCCESS; | |
| } | |
| int32_t GenericPluginTy::load_binary(int32_t DeviceId, | |
| __tgt_device_image *TgtImage, | |
| __tgt_device_binary *Binary) { | |
| GenericDeviceTy &Device = getDevice(DeviceId); | |
| StringRef Buffer(reinterpret_cast<const char *>(TgtImage->ImageStart), | |
| utils::getPtrDiff(TgtImage->ImageEnd, TgtImage->ImageStart)); | |
| auto ImageOrErr = Device.loadBinary(*this, Buffer); | |
| if (!ImageOrErr) { | |
| auto Err = ImageOrErr.takeError(); | |
| REPORT() << "Failure to load binary image " << TgtImage << " on device " | |
| << DeviceId << ": " << toString(std::move(Err)); | |
| return OFFLOAD_FAIL; | |
| } | |
| DeviceImageTy *Image = *ImageOrErr; | |
| assert(Image != nullptr && "Invalid Image"); | |
| *Binary = __tgt_device_binary{reinterpret_cast<uint64_t>(Image)}; | |
| return OFFLOAD_SUCCESS; | |
| } | |
| void *GenericPluginTy::data_alloc(int32_t DeviceId, int64_t Size, void *HostPtr, | |
| int32_t Kind) { | |
| auto AllocOrErr = getDevice(DeviceId).dataAlloc( | |
| Size, HostPtr, (TargetAllocTy)Kind, /*Alignment=*/0); | |
| if (!AllocOrErr) { | |
| auto Err = AllocOrErr.takeError(); | |
| REPORT() << "Failure to allocate device memory: " | |
| << toString(std::move(Err)); | |
| return nullptr; | |
| } | |
| assert(*AllocOrErr && "Null pointer upon successful allocation"); | |
| return *AllocOrErr; | |
| } | |
| int32_t GenericPluginTy::data_delete(int32_t DeviceId, void *TgtPtr, | |
| int32_t Kind) { | |
| auto Err = | |
| getDevice(DeviceId).dataDelete(TgtPtr, static_cast<TargetAllocTy>(Kind)); | |
| if (Err) { | |
| REPORT() << "Failure to deallocate device pointer " << TgtPtr << ": " | |
| << toString(std::move(Err)); | |
| return OFFLOAD_FAIL; | |
| } | |
| return OFFLOAD_SUCCESS; | |
| } | |
| int32_t GenericPluginTy::data_lock(int32_t DeviceId, void *Ptr, int64_t Size, | |
| void **LockedPtr) { | |
| auto LockedPtrOrErr = getDevice(DeviceId).registerMemory(Ptr, Size); | |
| if (!LockedPtrOrErr) { | |
| auto Err = LockedPtrOrErr.takeError(); | |
| REPORT() << "Failure to lock memory " << Ptr << ": " | |
| << toString(std::move(Err)); | |
| return OFFLOAD_FAIL; | |
| } | |
| if (!(*LockedPtrOrErr)) { | |
| REPORT() << "Failure to lock memory " << Ptr | |
| << ": obtained a null locked pointer"; | |
| return OFFLOAD_FAIL; | |
| } | |
| *LockedPtr = *LockedPtrOrErr; | |
| return OFFLOAD_SUCCESS; | |
| } | |
| int32_t GenericPluginTy::data_unlock(int32_t DeviceId, void *Ptr) { | |
| auto Err = getDevice(DeviceId).unregisterMemory(Ptr); | |
| if (Err) { | |
| REPORT() << "Failure to unlock memory " << Ptr << ": " | |
| << toString(std::move(Err)); | |
| return OFFLOAD_FAIL; | |
| } | |
| return OFFLOAD_SUCCESS; | |
| } | |
| int32_t GenericPluginTy::data_notify_mapped(int32_t DeviceId, void *HstPtr, | |
| int64_t Size) { | |
| auto Err = getDevice(DeviceId).notifyDataMapped(HstPtr, Size); | |
| if (Err) { | |
| REPORT() << "Failure to notify data mapped " << HstPtr << ": " | |
| << toString(std::move(Err)); | |
| return OFFLOAD_FAIL; | |
| } | |
| return OFFLOAD_SUCCESS; | |
| } | |
| int32_t GenericPluginTy::data_notify_unmapped(int32_t DeviceId, void *HstPtr) { | |
| auto Err = getDevice(DeviceId).notifyDataUnmapped(HstPtr); | |
| if (Err) { | |
| REPORT() << "Failure to notify data unmapped " << HstPtr << ": " | |
| << toString(std::move(Err)); | |
| return OFFLOAD_FAIL; | |
| } | |
| return OFFLOAD_SUCCESS; | |
| } | |
| int32_t GenericPluginTy::data_submit(int32_t DeviceId, void *TgtPtr, | |
| void *HstPtr, int64_t Size) { | |
| return data_submit_async(DeviceId, TgtPtr, HstPtr, Size, | |
| /*AsyncInfoPtr=*/nullptr); | |
| } | |
| int32_t GenericPluginTy::data_submit_async(int32_t DeviceId, void *TgtPtr, | |
| void *HstPtr, int64_t Size, | |
| __tgt_async_info *AsyncInfoPtr) { | |
| auto Err = getDevice(DeviceId).dataSubmit(TgtPtr, HstPtr, Size, AsyncInfoPtr); | |
| if (Err) { | |
| REPORT() << "Failure to copy data from host to device. Pointers: host " | |
| << "= " << HstPtr << ", device = " << TgtPtr << ", size = " << Size | |
| << ": " << toString(std::move(Err)); | |
| return OFFLOAD_FAIL; | |
| } | |
| return OFFLOAD_SUCCESS; | |
| } | |
| int32_t GenericPluginTy::data_retrieve(int32_t DeviceId, void *HstPtr, | |
| void *TgtPtr, int64_t Size) { | |
| return data_retrieve_async(DeviceId, HstPtr, TgtPtr, Size, | |
| /*AsyncInfoPtr=*/nullptr); | |
| } | |
| int32_t GenericPluginTy::data_retrieve_async(int32_t DeviceId, void *HstPtr, | |
| void *TgtPtr, int64_t Size, | |
| __tgt_async_info *AsyncInfoPtr) { | |
| auto Err = | |
| getDevice(DeviceId).dataRetrieve(HstPtr, TgtPtr, Size, AsyncInfoPtr); | |
| if (Err) { | |
| REPORT() << "Failure to copy data from device to host. Pointers: host " | |
| << "= " << HstPtr << ", device = " << TgtPtr << ", size = " << Size | |
| << ": " << toString(std::move(Err)); | |
| return OFFLOAD_FAIL; | |
| } | |
| return OFFLOAD_SUCCESS; | |
| } | |
| int32_t GenericPluginTy::data_exchange(int32_t SrcDeviceId, void *SrcPtr, | |
| int32_t DstDeviceId, void *DstPtr, | |
| int64_t Size) { | |
| return data_exchange_async(SrcDeviceId, SrcPtr, DstDeviceId, DstPtr, Size, | |
| /*AsyncInfoPtr=*/nullptr); | |
| } | |
| int32_t GenericPluginTy::data_exchange_async(int32_t SrcDeviceId, void *SrcPtr, | |
| int DstDeviceId, void *DstPtr, | |
| int64_t Size, | |
| __tgt_async_info *AsyncInfo) { | |
| GenericDeviceTy &SrcDevice = getDevice(SrcDeviceId); | |
| GenericDeviceTy &DstDevice = getDevice(DstDeviceId); | |
| auto Err = SrcDevice.dataExchange(SrcPtr, DstDevice, DstPtr, Size, AsyncInfo); | |
| if (Err) { | |
| REPORT() << "Failure to copy data from device (" << SrcDeviceId | |
| << ") to device (" << DstDeviceId | |
| << "). Pointers: host = " << SrcPtr << ", device = " << DstPtr | |
| << ", size = " << Size << ": " << toString(std::move(Err)); | |
| return OFFLOAD_FAIL; | |
| } | |
| return OFFLOAD_SUCCESS; | |
| } | |
| int32_t GenericPluginTy::launch_kernel(int32_t DeviceId, void *TgtEntryPtr, | |
| void **TgtArgs, ptrdiff_t *TgtOffsets, | |
| KernelArgsTy *KernelArgs, | |
| KernelExtraArgsTy *KernelExtraArgs, | |
| __tgt_async_info *AsyncInfoPtr) { | |
| auto Err = getDevice(DeviceId).launchKernel(TgtEntryPtr, TgtArgs, TgtOffsets, | |
| *KernelArgs, KernelExtraArgs, | |
| AsyncInfoPtr); | |
| if (Err) { | |
| REPORT() << "Failure to run target region " << TgtEntryPtr << " in device " | |
| << DeviceId << ": " << toString(std::move(Err)); | |
| return OFFLOAD_FAIL; | |
| } | |
| return OFFLOAD_SUCCESS; | |
| } | |
| int32_t GenericPluginTy::synchronize(int32_t DeviceId, | |
| __tgt_async_info *AsyncInfoPtr) { | |
| auto Err = getDevice(DeviceId).synchronize(AsyncInfoPtr); | |
| if (Err) { | |
| REPORT() << "Failure to synchronize stream " << AsyncInfoPtr->Queue << ": " | |
| << toString(std::move(Err)); | |
| return OFFLOAD_FAIL; | |
| } | |
| return OFFLOAD_SUCCESS; | |
| } | |
| int32_t GenericPluginTy::query_async(int32_t DeviceId, | |
| __tgt_async_info *AsyncInfoPtr) { | |
| auto Err = getDevice(DeviceId).queryAsync(AsyncInfoPtr); | |
| if (Err) { | |
| REPORT() << "Failure to query stream " << AsyncInfoPtr->Queue << ": " | |
| << toString(std::move(Err)); | |
| return OFFLOAD_FAIL; | |
| } | |
| return OFFLOAD_SUCCESS; | |
| } | |
| InfoTreeNode GenericPluginTy::obtain_device_info(int32_t DeviceId) { | |
| auto InfoOrErr = getDevice(DeviceId).obtainInfo(); | |
| if (auto Err = InfoOrErr.takeError()) { | |
| REPORT() << "Failure to obtain device " << DeviceId | |
| << " info: " << toString(std::move(Err)); | |
| return InfoTreeNode{}; | |
| } | |
| return std::move(*InfoOrErr); | |
| } | |
| void GenericPluginTy::print_device_info(int32_t DeviceId) { | |
| if (auto Err = getDevice(DeviceId).printInfo()) | |
| REPORT() << "Failure to print device " << DeviceId | |
| << " info: " << toString(std::move(Err)); | |
| } | |
| int32_t GenericPluginTy::create_event(int32_t DeviceId, void **EventPtr) { | |
| auto Err = getDevice(DeviceId).createEvent(EventPtr); | |
| if (Err) { | |
| REPORT() << "Failure to create event: " << toString(std::move(Err)); | |
| return OFFLOAD_FAIL; | |
| } | |
| return OFFLOAD_SUCCESS; | |
| } | |
| int32_t GenericPluginTy::record_event(int32_t DeviceId, void *EventPtr, | |
| __tgt_async_info *AsyncInfoPtr) { | |
| auto Err = getDevice(DeviceId).recordEvent(EventPtr, AsyncInfoPtr); | |
| if (Err) { | |
| REPORT() << "Failure to record event " << EventPtr << ": " | |
| << toString(std::move(Err)); | |
| return OFFLOAD_FAIL; | |
| } | |
| return OFFLOAD_SUCCESS; | |
| } | |
| int32_t GenericPluginTy::wait_event(int32_t DeviceId, void *EventPtr, | |
| __tgt_async_info *AsyncInfoPtr) { | |
| auto Err = getDevice(DeviceId).waitEvent(EventPtr, AsyncInfoPtr); | |
| if (Err) { | |
| REPORT() << "Failure to wait event " << EventPtr << ": " | |
| << toString(std::move(Err)); | |
| return OFFLOAD_FAIL; | |
| } | |
| return OFFLOAD_SUCCESS; | |
| } | |
| int32_t GenericPluginTy::sync_event(int32_t DeviceId, void *EventPtr) { | |
| auto Err = getDevice(DeviceId).syncEvent(EventPtr); | |
| if (Err) { | |
| REPORT() << "Failure to synchronize event " << EventPtr << ": " | |
| << toString(std::move(Err)); | |
| return OFFLOAD_FAIL; | |
| } | |
| return OFFLOAD_SUCCESS; | |
| } | |
| int32_t GenericPluginTy::get_event_elapsed_time(int32_t DeviceId, | |
| void *StartEventPtr, | |
| void *EndEventPtr, | |
| float *ElapsedTime) { | |
| auto ElapsedTimeOrErr = | |
| getDevice(DeviceId).getEventElapsedTime(StartEventPtr, EndEventPtr); | |
| if (!ElapsedTimeOrErr) { | |
| REPORT() << "Failure to get elapsed time between events " << StartEventPtr | |
| << " and " << EndEventPtr << ": " | |
| << toString(ElapsedTimeOrErr.takeError()); | |
| return OFFLOAD_FAIL; | |
| } | |
| *ElapsedTime = *ElapsedTimeOrErr; | |
| return OFFLOAD_SUCCESS; | |
| } | |
| int32_t GenericPluginTy::destroy_event(int32_t DeviceId, void *EventPtr) { | |
| auto Err = getDevice(DeviceId).destroyEvent(EventPtr); | |
| if (Err) { | |
| REPORT() << "Failure to destroy event " << EventPtr << ": " | |
| << toString(std::move(Err)); | |
| return OFFLOAD_FAIL; | |
| } | |
| return OFFLOAD_SUCCESS; | |
| } | |
| void GenericPluginTy::set_info_flag(uint32_t NewInfoLevel) { | |
| std::atomic<uint32_t> &InfoLevel = getInfoLevelInternal(); | |
| InfoLevel.store(NewInfoLevel); | |
| } | |
| int32_t GenericPluginTy::init_async_info(int32_t DeviceId, | |
| __tgt_async_info **AsyncInfoPtr) { | |
| assert(AsyncInfoPtr && "Invalid async info"); | |
| auto Err = getDevice(DeviceId).initAsyncInfo(AsyncInfoPtr); | |
| if (Err) { | |
| REPORT() << "Failure to initialize async info at " << *AsyncInfoPtr | |
| << " on device " << DeviceId << ": " << toString(std::move(Err)); | |
| return OFFLOAD_FAIL; | |
| } | |
| return OFFLOAD_SUCCESS; | |
| } | |
| int32_t GenericPluginTy::set_device_identifier(int32_t UserId, | |
| int32_t DeviceId) { | |
| UserDeviceIds[DeviceId] = UserId; | |
| return OFFLOAD_SUCCESS; | |
| } | |
| int32_t GenericPluginTy::use_auto_zero_copy(int32_t DeviceId) { | |
| return getDevice(DeviceId).useAutoZeroCopy(); | |
| } | |
| int32_t GenericPluginTy::is_accessible_ptr(int32_t DeviceId, const void *Ptr, | |
| size_t Size) { | |
| auto HandleError = [&](Error Err) -> bool { | |
| std::string ErrStr = toString(std::move(Err)); | |
| ODBG(OLDT_Device) << "Failure while checking accessibility of pointer " | |
| << Ptr << " for device " << DeviceId << ": " << ErrStr; | |
| return false; | |
| }; | |
| auto AccessibleOrErr = getDevice(DeviceId).isAccessiblePtr(Ptr, Size); | |
| if (Error Err = AccessibleOrErr.takeError()) | |
| return HandleError(std::move(Err)); | |
| return *AccessibleOrErr; | |
| } | |
| int32_t GenericPluginTy::get_global(__tgt_device_binary Binary, uint64_t Size, | |
| const char *Name, void **DevicePtr) { | |
| assert(Binary.handle && "Invalid device binary handle"); | |
| DeviceImageTy &Image = *reinterpret_cast<DeviceImageTy *>(Binary.handle); | |
| GenericDeviceTy &Device = Image.getDevice(); | |
| GlobalTy DeviceGlobal(Name, Size); | |
| GenericGlobalHandlerTy &GHandler = getGlobalHandler(); | |
| if (auto Err = | |
| GHandler.getGlobalMetadataFromDevice(Device, Image, DeviceGlobal)) { | |
| consumeError(std::move(Err)); | |
| return OFFLOAD_FAIL; | |
| } | |
| *DevicePtr = DeviceGlobal.getPtr(); | |
| assert(DevicePtr && "Invalid device global's address"); | |
| // Save the loaded globals if we are recording. | |
| RecordReplayTy *RecordReplay = Device.getRecordReplay(); | |
| if (RecordReplay && RecordReplay->isRecording()) | |
| RecordReplay->addGlobal(Name, Size, *DevicePtr); | |
| return OFFLOAD_SUCCESS; | |
| } | |
| int32_t GenericPluginTy::get_function(__tgt_device_binary Binary, | |
| const char *Name, void **KernelPtr) { | |
| assert(Binary.handle && "Invalid device binary handle"); | |
| DeviceImageTy &Image = *reinterpret_cast<DeviceImageTy *>(Binary.handle); | |
| GenericDeviceTy &Device = Image.getDevice(); | |
| auto KernelOrErr = Device.constructKernel(Name); | |
| if (Error Err = KernelOrErr.takeError()) { | |
| REPORT() << "Failure to look up kernel: " << toString(std::move(Err)); | |
| return OFFLOAD_FAIL; | |
| } | |
| GenericKernelTy &Kernel = *KernelOrErr; | |
| if (auto Err = Kernel.init(Device, Image)) { | |
| REPORT() << "Failure to init kernel: " << toString(std::move(Err)); | |
| return OFFLOAD_FAIL; | |
| } | |
| // Note that this is not the kernel's device address. | |
| *KernelPtr = &Kernel; | |
| return OFFLOAD_SUCCESS; | |
| } | |
| /// Create OpenMP interop with the given interop context | |
| omp_interop_val_t * | |
| GenericPluginTy::create_interop(int32_t ID, int32_t InteropContext, | |
| interop_spec_t *InteropSpec) { | |
| assert(InteropSpec && "Interop spec is null"); | |
| auto &Device = getDevice(ID); | |
| auto InteropOrErr = Device.createInterop(InteropContext, *InteropSpec); | |
| if (!InteropOrErr) { | |
| REPORT() << "Failure to create interop object for device " << InteropSpec | |
| << ": " << toString(InteropOrErr.takeError()); | |
| return nullptr; | |
| } | |
| return *InteropOrErr; | |
| } | |
| /// Release OpenMP interop object | |
| int32_t GenericPluginTy::release_interop(int32_t ID, | |
| omp_interop_val_t *Interop) { | |
| assert(Interop && "Interop is null"); | |
| assert(Interop->device_id == ID && "Interop does not match device id"); | |
| auto &Device = getDevice(ID); | |
| auto Err = Device.releaseInterop(Interop); | |
| if (Err) { | |
| REPORT() << "Failure to release interop object " << Interop << ": " | |
| << toString(std::move(Err)); | |
| return OFFLOAD_FAIL; | |
| } | |
| return OFFLOAD_SUCCESS; | |
| } | |
| /// Flush the queue associated with the interop object if necessary | |
| int32_t GenericPluginTy::flush_queue(omp_interop_val_t *Interop) { | |
| assert(Interop && "Interop is null"); | |
| auto Err = flushQueueImpl(Interop); | |
| if (Err) { | |
| REPORT() << "Failure to flush interop object " << Interop | |
| << " queue: " << toString(std::move(Err)); | |
| return OFFLOAD_FAIL; | |
| } | |
| return OFFLOAD_SUCCESS; | |
| } | |
| /// Perform a host synchronization with the queue associated with the interop | |
| /// object and wait for it to complete. | |
| int32_t GenericPluginTy::sync_barrier(omp_interop_val_t *Interop) { | |
| assert(Interop && "Interop is null"); | |
| auto Err = syncBarrierImpl(Interop); | |
| if (Err) { | |
| REPORT() << "Failure to synchronize interop object " << Interop << ": " | |
| << toString(std::move(Err)); | |
| return OFFLOAD_FAIL; | |
| } | |
| return OFFLOAD_SUCCESS; | |
| } | |
| /// Queue an asynchronous barrier in the queue associated with the interop | |
| /// object and return immediately. | |
| int32_t GenericPluginTy::async_barrier(omp_interop_val_t *Interop) { | |
| assert(Interop && "Interop is null"); | |
| auto Err = asyncBarrierImpl(Interop); | |
| if (Err) { | |
| REPORT() << "Failure to queue barrier in interop object " << Interop << ": " | |
| << toString(std::move(Err)); | |
| return OFFLOAD_FAIL; | |
| } | |
| return OFFLOAD_SUCCESS; | |
| } | |
| int32_t GenericPluginTy::data_fence(int32_t DeviceId, | |
| __tgt_async_info *AsyncInfo) { | |
| auto Err = getDevice(DeviceId).dataFence(AsyncInfo); | |
| if (Err) { | |
| REPORT() << "Failure to place data fence on device " << DeviceId << ": " | |
| << toString(std::move(Err)); | |
| return OFFLOAD_FAIL; | |
| } | |
| return OFFLOAD_SUCCESS; | |
| } |