1717 lines · cpp
1//===----RTLs/cuda/src/rtl.cpp - Target RTLs Implementation ------- C++ -*-===//2//3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.4// See https://llvm.org/LICENSE.txt for license information.5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception6//7//===----------------------------------------------------------------------===//8//9// RTL NextGen for CUDA machine10//11//===----------------------------------------------------------------------===//12 13#include <cassert>14#include <cstddef>15#include <cuda.h>16#include <string>17#include <unordered_map>18 19#include "Shared/APITypes.h"20#include "Shared/Debug.h"21#include "Shared/Environment.h"22 23#include "GlobalHandler.h"24#include "OpenMP/OMPT/Callback.h"25#include "PluginInterface.h"26#include "Utils/ELF.h"27 28#include "llvm/ADT/StringExtras.h"29#include "llvm/BinaryFormat/ELF.h"30#include "llvm/Frontend/OpenMP/OMPConstants.h"31#include "llvm/Frontend/OpenMP/OMPGridValues.h"32#include "llvm/Support/Error.h"33#include "llvm/Support/FileOutputBuffer.h"34#include "llvm/Support/FileSystem.h"35#include "llvm/Support/Program.h"36 37using namespace error;38 39namespace llvm {40namespace omp {41namespace target {42namespace plugin {43 44/// Forward declarations for all specialized data structures.45struct CUDAKernelTy;46struct CUDADeviceTy;47struct CUDAPluginTy;48 49#if (defined(CUDA_VERSION) && (CUDA_VERSION < 11000))50/// Forward declarations for all Virtual Memory Management51/// related data structures and functions. This is necessary52/// for older cuda versions.53typedef void *CUmemGenericAllocationHandle;54typedef void *CUmemAllocationProp;55typedef void *CUmemAccessDesc;56typedef void *CUmemAllocationGranularity_flags;57CUresult cuMemAddressReserve(CUdeviceptr *ptr, size_t size, size_t alignment,58 CUdeviceptr addr, unsigned long long flags) {}59CUresult cuMemMap(CUdeviceptr ptr, size_t size, size_t offset,60 CUmemGenericAllocationHandle handle,61 unsigned long long flags) {}62CUresult cuMemCreate(CUmemGenericAllocationHandle *handle, size_t size,63 const CUmemAllocationProp *prop,64 unsigned long long flags) {}65CUresult cuMemSetAccess(CUdeviceptr ptr, size_t size,66 const CUmemAccessDesc *desc, size_t count) {}67CUresult68cuMemGetAllocationGranularity(size_t *granularity,69 const CUmemAllocationProp *prop,70 CUmemAllocationGranularity_flags option) {}71#endif72 73#if (defined(CUDA_VERSION) && (CUDA_VERSION < 11020))74// Forward declarations of asynchronous memory management functions. This is75// necessary for older versions of CUDA.76CUresult cuMemAllocAsync(CUdeviceptr *ptr, size_t, CUstream) { *ptr = 0; }77 78CUresult cuMemFreeAsync(CUdeviceptr dptr, CUstream hStream) {}79#endif80 81/// Class implementing the CUDA device images properties.82struct CUDADeviceImageTy : public DeviceImageTy {83 /// Create the CUDA image with the id and the target image pointer.84 CUDADeviceImageTy(int32_t ImageId, GenericDeviceTy &Device,85 std::unique_ptr<MemoryBuffer> &&TgtImage)86 : DeviceImageTy(ImageId, Device, std::move(TgtImage)), Module(nullptr) {}87 88 /// Load the image as a CUDA module.89 Error loadModule() {90 assert(!Module && "Module already loaded");91 92 CUresult Res = cuModuleLoadDataEx(&Module, getStart(), 0, nullptr, nullptr);93 if (auto Err = Plugin::check(Res, "error in cuModuleLoadDataEx: %s"))94 return Err;95 96 return Plugin::success();97 }98 99 /// Unload the CUDA module corresponding to the image.100 Error unloadModule() {101 assert(Module && "Module not loaded");102 103 CUresult Res = cuModuleUnload(Module);104 if (auto Err = Plugin::check(Res, "error in cuModuleUnload: %s"))105 return Err;106 107 Module = nullptr;108 109 return Plugin::success();110 }111 112 /// Getter of the CUDA module.113 CUmodule getModule() const { return Module; }114 115private:116 /// The CUDA module that loaded the image.117 CUmodule Module;118};119 120/// Class implementing the CUDA kernel functionalities which derives from the121/// generic kernel class.122struct CUDAKernelTy : public GenericKernelTy {123 /// Create a CUDA kernel with a name and an execution mode.124 CUDAKernelTy(const char *Name) : GenericKernelTy(Name), Func(nullptr) {}125 126 /// Initialize the CUDA kernel.127 Error initImpl(GenericDeviceTy &GenericDevice,128 DeviceImageTy &Image) override {129 CUresult Res;130 CUDADeviceImageTy &CUDAImage = static_cast<CUDADeviceImageTy &>(Image);131 132 // Retrieve the function pointer of the kernel.133 Res = cuModuleGetFunction(&Func, CUDAImage.getModule(), getName());134 if (auto Err = Plugin::check(Res, "error in cuModuleGetFunction('%s'): %s",135 getName()))136 return Err;137 138 // Check that the function pointer is valid.139 if (!Func)140 return Plugin::error(ErrorCode::INVALID_BINARY,141 "invalid function for kernel %s", getName());142 143 int MaxThreads;144 Res = cuFuncGetAttribute(&MaxThreads,145 CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK, Func);146 if (auto Err = Plugin::check(Res, "error in cuFuncGetAttribute: %s"))147 return Err;148 149 // The maximum number of threads cannot exceed the maximum of the kernel.150 MaxNumThreads = std::min(MaxNumThreads, (uint32_t)MaxThreads);151 152 return Plugin::success();153 }154 155 /// Launch the CUDA kernel function.156 Error launchImpl(GenericDeviceTy &GenericDevice, uint32_t NumThreads[3],157 uint32_t NumBlocks[3], KernelArgsTy &KernelArgs,158 KernelLaunchParamsTy LaunchParams,159 AsyncInfoWrapperTy &AsyncInfoWrapper) const override;160 161 /// Return maximum block size for maximum occupancy162 Expected<uint64_t> maxGroupSize(GenericDeviceTy &,163 uint64_t DynamicMemSize) const override {164 int MinGridSize;165 int MaxBlockSize;166 auto Res = cuOccupancyMaxPotentialBlockSize(167 &MinGridSize, &MaxBlockSize, Func, NULL, DynamicMemSize, INT_MAX);168 if (auto Err = Plugin::check(169 Res, "error in cuOccupancyMaxPotentialBlockSize: %s")) {170 return Err;171 }172 return MaxBlockSize;173 }174 175private:176 /// The CUDA kernel function to execute.177 CUfunction Func;178 /// The maximum amount of dynamic shared memory per thread group. By default,179 /// this is set to 48 KB.180 mutable uint32_t MaxDynCGroupMemLimit = 49152;181};182 183/// Class wrapping a CUDA stream reference. These are the objects handled by the184/// Stream Manager for the CUDA plugin.185struct CUDAStreamRef final : public GenericDeviceResourceRef {186 /// The underlying handle type for streams.187 using HandleTy = CUstream;188 189 /// Create an empty reference to an invalid stream.190 CUDAStreamRef() : Stream(nullptr) {}191 192 /// Create a reference to an existing stream.193 CUDAStreamRef(HandleTy Stream) : Stream(Stream) {}194 195 /// Create a new stream and save the reference. The reference must be empty196 /// before calling to this function.197 Error create(GenericDeviceTy &Device) override {198 if (Stream)199 return Plugin::error(ErrorCode::INVALID_ARGUMENT,200 "creating an existing stream");201 202 CUresult Res = cuStreamCreate(&Stream, CU_STREAM_NON_BLOCKING);203 if (auto Err = Plugin::check(Res, "error in cuStreamCreate: %s"))204 return Err;205 206 return Plugin::success();207 }208 209 /// Destroy the referenced stream and invalidate the reference. The reference210 /// must be to a valid stream before calling to this function.211 Error destroy(GenericDeviceTy &Device) override {212 if (!Stream)213 return Plugin::error(ErrorCode::INVALID_ARGUMENT,214 "destroying an invalid stream");215 216 CUresult Res = cuStreamDestroy(Stream);217 if (auto Err = Plugin::check(Res, "error in cuStreamDestroy: %s"))218 return Err;219 220 Stream = nullptr;221 return Plugin::success();222 }223 224 /// Get the underlying CUDA stream.225 operator HandleTy() const { return Stream; }226 227private:228 /// The reference to the CUDA stream.229 HandleTy Stream;230};231 232/// Class wrapping a CUDA event reference. These are the objects handled by the233/// Event Manager for the CUDA plugin.234struct CUDAEventRef final : public GenericDeviceResourceRef {235 /// The underlying handle type for events.236 using HandleTy = CUevent;237 238 /// Create an empty reference to an invalid event.239 CUDAEventRef() : Event(nullptr) {}240 241 /// Create a reference to an existing event.242 CUDAEventRef(HandleTy Event) : Event(Event) {}243 244 /// Create a new event and save the reference. The reference must be empty245 /// before calling to this function.246 Error create(GenericDeviceTy &Device) override {247 if (Event)248 return Plugin::error(ErrorCode::INVALID_ARGUMENT,249 "creating an existing event");250 251 CUresult Res = cuEventCreate(&Event, CU_EVENT_DEFAULT);252 if (auto Err = Plugin::check(Res, "error in cuEventCreate: %s"))253 return Err;254 255 return Plugin::success();256 }257 258 /// Destroy the referenced event and invalidate the reference. The reference259 /// must be to a valid event before calling to this function.260 Error destroy(GenericDeviceTy &Device) override {261 if (!Event)262 return Plugin::error(ErrorCode::INVALID_ARGUMENT,263 "destroying an invalid event");264 265 CUresult Res = cuEventDestroy(Event);266 if (auto Err = Plugin::check(Res, "error in cuEventDestroy: %s"))267 return Err;268 269 Event = nullptr;270 return Plugin::success();271 }272 273 /// Get the underlying CUevent.274 operator HandleTy() const { return Event; }275 276private:277 /// The reference to the CUDA event.278 HandleTy Event;279};280 281/// Class implementing the CUDA device functionalities which derives from the282/// generic device class.283struct CUDADeviceTy : public GenericDeviceTy {284 // Create a CUDA device with a device id and the default CUDA grid values.285 CUDADeviceTy(GenericPluginTy &Plugin, int32_t DeviceId, int32_t NumDevices)286 : GenericDeviceTy(Plugin, DeviceId, NumDevices, NVPTXGridValues),287 CUDAStreamManager(*this), CUDAEventManager(*this) {}288 289 ~CUDADeviceTy() {}290 291 /// Initialize the device, its resources and get its properties.292 Error initImpl(GenericPluginTy &Plugin) override {293 CUresult Res = cuDeviceGet(&Device, DeviceId);294 if (auto Err = Plugin::check(Res, "error in cuDeviceGet: %s"))295 return Err;296 297 CUuuid UUID = {0};298 Res = cuDeviceGetUuid(&UUID, Device);299 if (auto Err = Plugin::check(Res, "error in cuDeviceGetUuid: %s"))300 return Err;301 setDeviceUidFromVendorUid(toHex(UUID.bytes, true));302 303 // Query the current flags of the primary context and set its flags if304 // it is inactive.305 unsigned int FormerPrimaryCtxFlags = 0;306 int FormerPrimaryCtxIsActive = 0;307 Res = cuDevicePrimaryCtxGetState(Device, &FormerPrimaryCtxFlags,308 &FormerPrimaryCtxIsActive);309 if (auto Err =310 Plugin::check(Res, "error in cuDevicePrimaryCtxGetState: %s"))311 return Err;312 313 if (FormerPrimaryCtxIsActive) {314 INFO(OMP_INFOTYPE_PLUGIN_KERNEL, DeviceId,315 "The primary context is active, no change to its flags\n");316 if ((FormerPrimaryCtxFlags & CU_CTX_SCHED_MASK) !=317 CU_CTX_SCHED_BLOCKING_SYNC)318 INFO(OMP_INFOTYPE_PLUGIN_KERNEL, DeviceId,319 "Warning: The current flags are not CU_CTX_SCHED_BLOCKING_SYNC\n");320 } else {321 INFO(OMP_INFOTYPE_PLUGIN_KERNEL, DeviceId,322 "The primary context is inactive, set its flags to "323 "CU_CTX_SCHED_BLOCKING_SYNC\n");324 Res = cuDevicePrimaryCtxSetFlags(Device, CU_CTX_SCHED_BLOCKING_SYNC);325 if (auto Err =326 Plugin::check(Res, "error in cuDevicePrimaryCtxSetFlags: %s"))327 return Err;328 }329 330 // Retain the per device primary context and save it to use whenever this331 // device is selected.332 Res = cuDevicePrimaryCtxRetain(&Context, Device);333 if (auto Err = Plugin::check(Res, "error in cuDevicePrimaryCtxRetain: %s"))334 return Err;335 336 if (auto Err = setContext())337 return Err;338 339 // Initialize stream pool.340 if (auto Err = CUDAStreamManager.init(OMPX_InitialNumStreams))341 return Err;342 343 // Initialize event pool.344 if (auto Err = CUDAEventManager.init(OMPX_InitialNumEvents))345 return Err;346 347 // Query attributes to determine number of threads/block and blocks/grid.348 if (auto Err = getDeviceAttr(CU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_X,349 GridValues.GV_Max_Teams))350 return Err;351 352 if (auto Err = getDeviceAttr(CU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_X,353 GridValues.GV_Max_WG_Size))354 return Err;355 356 if (auto Err = getDeviceAttr(CU_DEVICE_ATTRIBUTE_WARP_SIZE,357 GridValues.GV_Warp_Size))358 return Err;359 360 if (auto Err = getDeviceAttr(CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR,361 ComputeCapability.Major))362 return Err;363 364 if (auto Err = getDeviceAttr(CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR,365 ComputeCapability.Minor))366 return Err;367 368 uint32_t NumMuliprocessors = 0;369 uint32_t MaxThreadsPerSM = 0;370 uint32_t WarpSize = 0;371 if (auto Err = getDeviceAttr(CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT,372 NumMuliprocessors))373 return Err;374 if (auto Err =375 getDeviceAttr(CU_DEVICE_ATTRIBUTE_MAX_THREADS_PER_MULTIPROCESSOR,376 MaxThreadsPerSM))377 return Err;378 if (auto Err = getDeviceAttr(CU_DEVICE_ATTRIBUTE_WARP_SIZE, WarpSize))379 return Err;380 HardwareParallelism = NumMuliprocessors * (MaxThreadsPerSM / WarpSize);381 382 uint32_t MaxSharedMem;383 if (auto Err = getDeviceAttr(384 CU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_BLOCK, MaxSharedMem))385 return Err;386 MaxBlockSharedMemSize = MaxSharedMem;387 388 return Plugin::success();389 }390 391 Error unloadBinaryImpl(DeviceImageTy *Image) override {392 assert(Context && "Invalid CUDA context");393 394 // Each image has its own module.395 CUDADeviceImageTy &CUDAImage = static_cast<CUDADeviceImageTy &>(*Image);396 397 // Unload the module of the image.398 if (auto Err = CUDAImage.unloadModule())399 return Err;400 401 // Destroy the associated memory and invalidate the object.402 Plugin.free(Image);403 return Plugin::success();404 }405 406 /// Deinitialize the device and release its resources.407 Error deinitImpl() override {408 if (Context) {409 if (auto Err = setContext())410 return Err;411 }412 413 // Deinitialize the stream manager.414 if (auto Err = CUDAStreamManager.deinit())415 return Err;416 417 if (auto Err = CUDAEventManager.deinit())418 return Err;419 420 if (Context) {421 CUresult Res = cuDevicePrimaryCtxRelease(Device);422 if (auto Err =423 Plugin::check(Res, "error in cuDevicePrimaryCtxRelease: %s"))424 return Err;425 }426 427 // Invalidate context and device references.428 Context = nullptr;429 Device = CU_DEVICE_INVALID;430 431 return Plugin::success();432 }433 434 virtual Error callGlobalConstructors(GenericPluginTy &Plugin,435 DeviceImageTy &Image) override {436 return callGlobalCtorDtorCommon(Plugin, Image, /*IsCtor=*/true);437 }438 439 virtual Error callGlobalDestructors(GenericPluginTy &Plugin,440 DeviceImageTy &Image) override {441 return callGlobalCtorDtorCommon(Plugin, Image, /*IsCtor=*/false);442 }443 444 Expected<std::unique_ptr<MemoryBuffer>>445 doJITPostProcessing(std::unique_ptr<MemoryBuffer> MB) const override {446 // TODO: We should be able to use the 'nvidia-ptxjitcompiler' interface to447 // avoid the call to 'ptxas'.448 SmallString<128> PTXInputFilePath;449 std::error_code EC = sys::fs::createTemporaryFile("nvptx-pre-link-jit", "s",450 PTXInputFilePath);451 if (EC)452 return Plugin::error(ErrorCode::HOST_IO,453 "failed to create temporary file for ptxas");454 455 // Write the file's contents to the output file.456 Expected<std::unique_ptr<FileOutputBuffer>> OutputOrErr =457 FileOutputBuffer::create(PTXInputFilePath, MB->getBuffer().size());458 if (!OutputOrErr)459 return OutputOrErr.takeError();460 std::unique_ptr<FileOutputBuffer> Output = std::move(*OutputOrErr);461 llvm::copy(MB->getBuffer(), Output->getBufferStart());462 if (Error E = Output->commit())463 return std::move(E);464 465 SmallString<128> PTXOutputFilePath;466 EC = sys::fs::createTemporaryFile("nvptx-post-link-jit", "cubin",467 PTXOutputFilePath);468 if (EC)469 return Plugin::error(ErrorCode::HOST_IO,470 "failed to create temporary file for ptxas");471 472 // Try to find `ptxas` in the path to compile the PTX to a binary.473 const auto ErrorOrPath = sys::findProgramByName("ptxas");474 if (!ErrorOrPath)475 return Plugin::error(ErrorCode::HOST_TOOL_NOT_FOUND,476 "failed to find 'ptxas' on the PATH.");477 478 std::string Arch = getComputeUnitKind();479 StringRef Args[] = {*ErrorOrPath,480 "-m64",481 "-O2",482 "--gpu-name",483 Arch,484 "--output-file",485 PTXOutputFilePath,486 PTXInputFilePath};487 488 std::string ErrMsg;489 if (sys::ExecuteAndWait(*ErrorOrPath, Args, std::nullopt, {}, 0, 0,490 &ErrMsg))491 return Plugin::error(ErrorCode::ASSEMBLE_FAILURE,492 "running 'ptxas' failed: %s\n", ErrMsg.c_str());493 494 auto BufferOrErr = MemoryBuffer::getFileOrSTDIN(PTXOutputFilePath.data());495 if (!BufferOrErr)496 return Plugin::error(ErrorCode::HOST_IO,497 "failed to open temporary file for ptxas");498 499 // Clean up the temporary files afterwards.500 if (sys::fs::remove(PTXOutputFilePath))501 return Plugin::error(ErrorCode::HOST_IO,502 "failed to remove temporary file for ptxas");503 if (sys::fs::remove(PTXInputFilePath))504 return Plugin::error(ErrorCode::HOST_IO,505 "failed to remove temporary file for ptxas");506 507 return std::move(*BufferOrErr);508 }509 510 /// Allocate and construct a CUDA kernel.511 Expected<GenericKernelTy &> constructKernel(const char *Name) override {512 // Allocate and construct the CUDA kernel.513 CUDAKernelTy *CUDAKernel = Plugin.allocate<CUDAKernelTy>();514 if (!CUDAKernel)515 return Plugin::error(ErrorCode::OUT_OF_RESOURCES,516 "failed to allocate memory for CUDA kernel");517 518 new (CUDAKernel) CUDAKernelTy(Name);519 520 return *CUDAKernel;521 }522 523 /// Set the current context to this device's context.524 Error setContext() override {525 CUresult Res = cuCtxSetCurrent(Context);526 return Plugin::check(Res, "error in cuCtxSetCurrent: %s");527 }528 529 /// NVIDIA returns the product of the SM count and the number of warps that530 /// fit if the maximum number of threads were scheduled on each SM.531 uint64_t getHardwareParallelism() const override {532 return HardwareParallelism;533 }534 535 /// We want to set up the RPC server for host services to the GPU if it is536 /// available.537 bool shouldSetupRPCServer() const override { return true; }538 539 /// The RPC interface should have enough space for all available parallelism.540 uint64_t requestedRPCPortCount() const override {541 return getHardwareParallelism();542 }543 544 /// Get the stream of the asynchronous info structure or get a new one.545 Error getStream(AsyncInfoWrapperTy &AsyncInfoWrapper, CUstream &Stream) {546 auto WrapperStream =547 AsyncInfoWrapper.getOrInitQueue<CUstream>(CUDAStreamManager);548 if (!WrapperStream)549 return WrapperStream.takeError();550 Stream = *WrapperStream;551 return Plugin::success();552 }553 554 /// Getters of CUDA references.555 CUcontext getCUDAContext() const { return Context; }556 CUdevice getCUDADevice() const { return Device; }557 558 /// Load the binary image into the device and allocate an image object.559 Expected<DeviceImageTy *>560 loadBinaryImpl(std::unique_ptr<MemoryBuffer> &&TgtImage,561 int32_t ImageId) override {562 if (auto Err = setContext())563 return std::move(Err);564 565 // Allocate and initialize the image object.566 CUDADeviceImageTy *CUDAImage = Plugin.allocate<CUDADeviceImageTy>();567 new (CUDAImage) CUDADeviceImageTy(ImageId, *this, std::move(TgtImage));568 569 // Load the CUDA module.570 if (auto Err = CUDAImage->loadModule())571 return std::move(Err);572 573 return CUDAImage;574 }575 576 /// Allocate memory on the device or related to the device.577 Expected<void *> allocate(size_t Size, void *, TargetAllocTy Kind) override {578 if (Size == 0)579 return nullptr;580 581 if (auto Err = setContext())582 return std::move(Err);583 584 void *MemAlloc = nullptr;585 CUdeviceptr DevicePtr;586 CUresult Res;587 588 switch (Kind) {589 case TARGET_ALLOC_DEFAULT:590 case TARGET_ALLOC_DEVICE:591 Res = cuMemAlloc(&DevicePtr, Size);592 MemAlloc = (void *)DevicePtr;593 break;594 case TARGET_ALLOC_HOST:595 Res = cuMemAllocHost(&MemAlloc, Size);596 break;597 case TARGET_ALLOC_SHARED:598 Res = cuMemAllocManaged(&DevicePtr, Size, CU_MEM_ATTACH_GLOBAL);599 MemAlloc = (void *)DevicePtr;600 break;601 }602 603 if (auto Err = Plugin::check(Res, "error in cuMemAlloc[Host|Managed]: %s"))604 return std::move(Err);605 return MemAlloc;606 }607 608 /// Deallocate memory on the device or related to the device.609 Error free(void *TgtPtr, TargetAllocTy Kind) override {610 if (TgtPtr == nullptr)611 return Plugin::success();612 613 if (auto Err = setContext())614 return Err;615 616 CUresult Res;617 switch (Kind) {618 case TARGET_ALLOC_DEFAULT:619 case TARGET_ALLOC_DEVICE:620 case TARGET_ALLOC_SHARED:621 Res = cuMemFree((CUdeviceptr)TgtPtr);622 break;623 case TARGET_ALLOC_HOST:624 Res = cuMemFreeHost(TgtPtr);625 break;626 }627 628 return Plugin::check(Res, "error in cuMemFree[Host]: %s");629 }630 631 /// Synchronize current thread with the pending operations on the async info.632 Error synchronizeImpl(__tgt_async_info &AsyncInfo,633 bool ReleaseQueue) override {634 CUstream Stream = reinterpret_cast<CUstream>(AsyncInfo.Queue);635 CUresult Res;636 Res = cuStreamSynchronize(Stream);637 638 // Once the stream is synchronized and we want to release the queue, return639 // it to stream pool and reset AsyncInfo. This is to make sure the640 // synchronization only works for its own tasks.641 if (ReleaseQueue) {642 AsyncInfo.Queue = nullptr;643 if (auto Err = CUDAStreamManager.returnResource(Stream))644 return Err;645 }646 647 return Plugin::check(Res, "error in cuStreamSynchronize: %s");648 }649 650 /// CUDA support VA management651 bool supportVAManagement() const override {652#if (defined(CUDA_VERSION) && (CUDA_VERSION >= 11000))653 return true;654#else655 return false;656#endif657 }658 659 /// Allocates \p RSize bytes (rounded up to page size) and hints the cuda660 /// driver to map it to \p VAddr. The obtained address is stored in \p Addr.661 /// At return \p RSize contains the actual size662 Error memoryVAMap(void **Addr, void *VAddr, size_t *RSize) override {663 CUdeviceptr DVAddr = reinterpret_cast<CUdeviceptr>(VAddr);664 auto IHandle = DeviceMMaps.find(DVAddr);665 size_t Size = *RSize;666 667 if (Size == 0)668 return Plugin::error(ErrorCode::INVALID_ARGUMENT,669 "memory Map Size must be larger than 0");670 671 // Check if we have already mapped this address672 if (IHandle != DeviceMMaps.end())673 return Plugin::error(ErrorCode::INVALID_ARGUMENT,674 "address already memory mapped");675 676 CUmemAllocationProp Prop = {};677 size_t Granularity = 0;678 679 size_t Free, Total;680 CUresult Res = cuMemGetInfo(&Free, &Total);681 if (auto Err = Plugin::check(Res, "Error in cuMemGetInfo: %s"))682 return Err;683 684 if (Size >= Free) {685 *Addr = nullptr;686 return Plugin::error(687 ErrorCode::OUT_OF_RESOURCES,688 "cannot map memory size larger than the available device memory");689 }690 691 // currently NVidia only supports pinned device types692 Prop.type = CU_MEM_ALLOCATION_TYPE_PINNED;693 Prop.location.type = CU_MEM_LOCATION_TYPE_DEVICE;694 695 Prop.location.id = DeviceId;696 cuMemGetAllocationGranularity(&Granularity, &Prop,697 CU_MEM_ALLOC_GRANULARITY_MINIMUM);698 if (auto Err =699 Plugin::check(Res, "error in cuMemGetAllocationGranularity: %s"))700 return Err;701 702 if (Granularity == 0)703 return Plugin::error(ErrorCode::INVALID_ARGUMENT,704 "wrong device Page size");705 706 // Ceil to page size.707 Size = utils::roundUp(Size, Granularity);708 709 // Create a handler of our allocation710 CUmemGenericAllocationHandle AHandle;711 Res = cuMemCreate(&AHandle, Size, &Prop, 0);712 if (auto Err = Plugin::check(Res, "error in cuMemCreate: %s"))713 return Err;714 715 CUdeviceptr DevPtr = 0;716 Res = cuMemAddressReserve(&DevPtr, Size, 0, DVAddr, 0);717 if (auto Err = Plugin::check(Res, "error in cuMemAddressReserve: %s"))718 return Err;719 720 Res = cuMemMap(DevPtr, Size, 0, AHandle, 0);721 if (auto Err = Plugin::check(Res, "error in cuMemMap: %s"))722 return Err;723 724 CUmemAccessDesc ADesc = {};725 ADesc.location.type = CU_MEM_LOCATION_TYPE_DEVICE;726 ADesc.location.id = DeviceId;727 ADesc.flags = CU_MEM_ACCESS_FLAGS_PROT_READWRITE;728 729 // Sets address730 Res = cuMemSetAccess(DevPtr, Size, &ADesc, 1);731 if (auto Err = Plugin::check(Res, "error in cuMemSetAccess: %s"))732 return Err;733 734 *Addr = reinterpret_cast<void *>(DevPtr);735 *RSize = Size;736 DeviceMMaps.insert({DevPtr, AHandle});737 return Plugin::success();738 }739 740 /// De-allocates device memory and Unmaps the Virtual Addr741 Error memoryVAUnMap(void *VAddr, size_t Size) override {742 CUdeviceptr DVAddr = reinterpret_cast<CUdeviceptr>(VAddr);743 auto IHandle = DeviceMMaps.find(DVAddr);744 // Mapping does not exist745 if (IHandle == DeviceMMaps.end()) {746 return Plugin::error(ErrorCode::INVALID_ARGUMENT,747 "addr is not MemoryMapped");748 }749 750 if (IHandle == DeviceMMaps.end())751 return Plugin::error(ErrorCode::INVALID_ARGUMENT,752 "addr is not MemoryMapped");753 754 CUmemGenericAllocationHandle &AllocHandle = IHandle->second;755 756 CUresult Res = cuMemUnmap(DVAddr, Size);757 if (auto Err = Plugin::check(Res, "error in cuMemUnmap: %s"))758 return Err;759 760 Res = cuMemRelease(AllocHandle);761 if (auto Err = Plugin::check(Res, "error in cuMemRelease: %s"))762 return Err;763 764 Res = cuMemAddressFree(DVAddr, Size);765 if (auto Err = Plugin::check(Res, "error in cuMemAddressFree: %s"))766 return Err;767 768 DeviceMMaps.erase(IHandle);769 return Plugin::success();770 }771 772 /// Query for the completion of the pending operations on the async info.773 Error queryAsyncImpl(__tgt_async_info &AsyncInfo) override {774 CUstream Stream = reinterpret_cast<CUstream>(AsyncInfo.Queue);775 CUresult Res = cuStreamQuery(Stream);776 777 // Not ready streams must be considered as successful operations.778 if (Res == CUDA_ERROR_NOT_READY)779 return Plugin::success();780 781 // Once the stream is synchronized and the operations completed (or an error782 // occurs), return it to stream pool and reset AsyncInfo. This is to make783 // sure the synchronization only works for its own tasks.784 AsyncInfo.Queue = nullptr;785 if (auto Err = CUDAStreamManager.returnResource(Stream))786 return Err;787 788 return Plugin::check(Res, "error in cuStreamQuery: %s");789 }790 791 Expected<void *> dataLockImpl(void *HstPtr, int64_t Size) override {792 // TODO: Register the buffer as CUDA host memory.793 return HstPtr;794 }795 796 Error dataUnlockImpl(void *HstPtr) override { return Plugin::success(); }797 798 Expected<bool> isPinnedPtrImpl(void *HstPtr, void *&BaseHstPtr,799 void *&BaseDevAccessiblePtr,800 size_t &BaseSize) const override {801 // TODO: Implement pinning feature for CUDA.802 return false;803 }804 805 /// Submit data to the device (host to device transfer).806 Error dataSubmitImpl(void *TgtPtr, const void *HstPtr, int64_t Size,807 AsyncInfoWrapperTy &AsyncInfoWrapper) override {808 if (auto Err = setContext())809 return Err;810 811 CUstream Stream;812 if (auto Err = getStream(AsyncInfoWrapper, Stream))813 return Err;814 815 CUresult Res = cuMemcpyHtoDAsync((CUdeviceptr)TgtPtr, HstPtr, Size, Stream);816 return Plugin::check(Res, "error in cuMemcpyHtoDAsync: %s");817 }818 819 /// Retrieve data from the device (device to host transfer).820 Error dataRetrieveImpl(void *HstPtr, const void *TgtPtr, int64_t Size,821 AsyncInfoWrapperTy &AsyncInfoWrapper) override {822 if (auto Err = setContext())823 return Err;824 825 CUstream Stream;826 if (auto Err = getStream(AsyncInfoWrapper, Stream))827 return Err;828 829 CUresult Res = cuMemcpyDtoHAsync(HstPtr, (CUdeviceptr)TgtPtr, Size, Stream);830 return Plugin::check(Res, "error in cuMemcpyDtoHAsync: %s");831 }832 833 /// Exchange data between two devices directly. We may use peer access if834 /// the CUDA devices and driver allow them.835 Error dataExchangeImpl(const void *SrcPtr, GenericDeviceTy &DstGenericDevice,836 void *DstPtr, int64_t Size,837 AsyncInfoWrapperTy &AsyncInfoWrapper) override;838 839 Error dataFillImpl(void *TgtPtr, const void *PatternPtr, int64_t PatternSize,840 int64_t Size,841 AsyncInfoWrapperTy &AsyncInfoWrapper) override {842 if (auto Err = setContext())843 return Err;844 845 CUstream Stream;846 if (auto Err = getStream(AsyncInfoWrapper, Stream))847 return Err;848 849 CUresult Res;850 size_t N = Size / PatternSize;851 if (PatternSize == 1) {852 Res = cuMemsetD8Async((CUdeviceptr)TgtPtr,853 *(static_cast<const uint8_t *>(PatternPtr)), N,854 Stream);855 } else if (PatternSize == 2) {856 Res = cuMemsetD16Async((CUdeviceptr)TgtPtr,857 *(static_cast<const uint16_t *>(PatternPtr)), N,858 Stream);859 } else if (PatternSize == 4) {860 Res = cuMemsetD32Async((CUdeviceptr)TgtPtr,861 *(static_cast<const uint32_t *>(PatternPtr)), N,862 Stream);863 } else {864 // For larger patterns we can do a series of strided fills to copy the865 // pattern efficiently866 int64_t MemsetSize = PatternSize % 4u == 0u ? 4u867 : PatternSize % 2u == 0u ? 2u868 : 1u;869 870 int64_t NumberOfSteps = PatternSize / MemsetSize;871 int64_t Pitch = NumberOfSteps * MemsetSize;872 int64_t Height = Size / PatternSize;873 874 for (auto Step = 0u; Step < NumberOfSteps; ++Step) {875 if (MemsetSize == 4) {876 Res = cuMemsetD2D32Async(877 (CUdeviceptr)TgtPtr + Step * MemsetSize, Pitch,878 *(static_cast<const uint32_t *>(PatternPtr) + Step), 1u, Height,879 Stream);880 } else if (MemsetSize == 2) {881 Res = cuMemsetD2D16Async(882 (CUdeviceptr)TgtPtr + Step * MemsetSize, Pitch,883 *(static_cast<const uint16_t *>(PatternPtr) + Step), 1u, Height,884 Stream);885 } else {886 Res = cuMemsetD2D8Async(887 (CUdeviceptr)TgtPtr + Step * MemsetSize, Pitch,888 *(static_cast<const uint8_t *>(PatternPtr) + Step), 1u, Height,889 Stream);890 }891 }892 }893 894 return Plugin::check(Res, "error in cuMemset: %s");895 }896 897 /// Initialize the async info for interoperability purposes.898 Error initAsyncInfoImpl(AsyncInfoWrapperTy &AsyncInfoWrapper) override {899 if (auto Err = setContext())900 return Err;901 902 CUstream Stream;903 if (auto Err = getStream(AsyncInfoWrapper, Stream))904 return Err;905 906 return Plugin::success();907 }908 909 /// Insert a data fence between previous data operations and the following910 /// operations. This is a no-op for CUDA devices as operations inserted into911 /// a queue are in-order.912 Error dataFence(__tgt_async_info *Async) override {913 return Plugin::success();914 }915 916 interop_spec_t selectInteropPreference(int32_t InteropType,917 int32_t NumPrefers,918 interop_spec_t *Prefers) override {919 return interop_spec_t{tgt_fr_cuda, {true, 0}, 0};920 }921 922 Expected<omp_interop_val_t *>923 createInterop(int32_t InteropType, interop_spec_t &InteropSpec) override {924 auto *Ret = new omp_interop_val_t(925 DeviceId, static_cast<kmp_interop_type_t>(InteropType));926 Ret->fr_id = tgt_fr_cuda;927 Ret->vendor_id = omp_vendor_nvidia;928 929 if (InteropType == kmp_interop_type_target ||930 InteropType == kmp_interop_type_targetsync) {931 Ret->device_info.Platform = nullptr;932 Ret->device_info.Device = reinterpret_cast<void *>(Device);933 Ret->device_info.Context = Context;934 }935 936 if (InteropType == kmp_interop_type_targetsync) {937 Ret->async_info = new __tgt_async_info();938 if (auto Err = setContext())939 return Err;940 CUstream Stream;941 if (auto Err = CUDAStreamManager.getResource(Stream))942 return Err;943 944 Ret->async_info->Queue = Stream;945 }946 return Ret;947 }948 949 Error releaseInterop(omp_interop_val_t *Interop) override {950 if (!Interop)951 return Plugin::success();952 953 if (Interop->async_info)954 delete Interop->async_info;955 956 delete Interop;957 return Plugin::success();958 }959 960 Error enqueueHostCallImpl(void (*Callback)(void *), void *UserData,961 AsyncInfoWrapperTy &AsyncInfo) override {962 if (auto Err = setContext())963 return Err;964 965 CUstream Stream;966 if (auto Err = getStream(AsyncInfo, Stream))967 return Err;968 969 CUresult Res = cuLaunchHostFunc(Stream, Callback, UserData);970 return Plugin::check(Res, "error in cuStreamLaunchHostFunc: %s");971 };972 973 /// Create an event.974 Error createEventImpl(void **EventPtrStorage) override {975 CUevent *Event = reinterpret_cast<CUevent *>(EventPtrStorage);976 return CUDAEventManager.getResource(*Event);977 }978 979 /// Destroy a previously created event.980 Error destroyEventImpl(void *EventPtr) override {981 CUevent Event = reinterpret_cast<CUevent>(EventPtr);982 return CUDAEventManager.returnResource(Event);983 }984 985 /// Record the event.986 Error recordEventImpl(void *EventPtr,987 AsyncInfoWrapperTy &AsyncInfoWrapper) override {988 CUevent Event = reinterpret_cast<CUevent>(EventPtr);989 990 CUstream Stream;991 if (auto Err = getStream(AsyncInfoWrapper, Stream))992 return Err;993 994 CUresult Res = cuEventRecord(Event, Stream);995 return Plugin::check(Res, "error in cuEventRecord: %s");996 }997 998 /// Make the stream wait on the event.999 Error waitEventImpl(void *EventPtr,1000 AsyncInfoWrapperTy &AsyncInfoWrapper) override {1001 CUevent Event = reinterpret_cast<CUevent>(EventPtr);1002 1003 CUstream Stream;1004 if (auto Err = getStream(AsyncInfoWrapper, Stream))1005 return Err;1006 1007 // Do not use CU_EVENT_WAIT_DEFAULT here as it is only available from1008 // specific CUDA version, and defined as 0x0. In previous version, per CUDA1009 // API document, that argument has to be 0x0.1010 CUresult Res = cuStreamWaitEvent(Stream, Event, 0);1011 return Plugin::check(Res, "error in cuStreamWaitEvent: %s");1012 }1013 1014 Expected<bool> hasPendingWorkImpl(AsyncInfoWrapperTy &AsyncInfo) override {1015 CUstream Stream;1016 if (auto Err = getStream(AsyncInfo, Stream))1017 return Err;1018 1019 CUresult Ret = cuStreamQuery(Stream);1020 if (Ret == CUDA_SUCCESS)1021 return false;1022 1023 if (Ret == CUDA_ERROR_NOT_READY)1024 return true;1025 1026 return Plugin::check(Ret, "error in cuStreamQuery: %s");1027 }1028 1029 Expected<bool> isEventCompleteImpl(void *EventPtr,1030 AsyncInfoWrapperTy &) override {1031 CUevent Event = reinterpret_cast<CUevent>(EventPtr);1032 1033 CUresult Ret = cuEventQuery(Event);1034 if (Ret == CUDA_SUCCESS)1035 return true;1036 1037 if (Ret == CUDA_ERROR_NOT_READY)1038 return false;1039 1040 return Plugin::check(Ret, "error in cuEventQuery: %s");1041 }1042 1043 /// Synchronize the current thread with the event.1044 Error syncEventImpl(void *EventPtr) override {1045 CUevent Event = reinterpret_cast<CUevent>(EventPtr);1046 CUresult Res = cuEventSynchronize(Event);1047 return Plugin::check(Res, "error in cuEventSynchronize: %s");1048 }1049 1050 /// Print information about the device.1051 Expected<InfoTreeNode> obtainInfoImpl() override {1052 char TmpChar[1000];1053 const char *TmpCharPtr;1054 size_t TmpSt;1055 int TmpInt;1056 InfoTreeNode Info;1057 1058 CUresult Res = cuDriverGetVersion(&TmpInt);1059 if (Res == CUDA_SUCCESS)1060 // For consistency with other drivers, store the version as a string1061 // rather than an integer1062 Info.add("CUDA Driver Version", std::to_string(TmpInt), "",1063 DeviceInfo::DRIVER_VERSION);1064 1065 Info.add("CUDA OpenMP Device Number", DeviceId);1066 1067 Res = cuDeviceGetName(TmpChar, 1000, Device);1068 if (Res == CUDA_SUCCESS) {1069 Info.add("Device Name", TmpChar, "", DeviceInfo::NAME);1070 Info.add("Product Name", TmpChar, "", DeviceInfo::PRODUCT_NAME);1071 }1072 1073 Info.add("Vendor Name", "NVIDIA", "", DeviceInfo::VENDOR);1074 1075 Info.add("Vendor ID", uint64_t{4318}, "", DeviceInfo::VENDOR_ID);1076 1077 Info.add("Memory Address Size", std::numeric_limits<CUdeviceptr>::digits,1078 "bits", DeviceInfo::ADDRESS_BITS);1079 1080 Res = cuDeviceTotalMem(&TmpSt, Device);1081 if (Res == CUDA_SUCCESS)1082 Info.add("Global Memory Size", TmpSt, "bytes",1083 DeviceInfo::GLOBAL_MEM_SIZE);1084 1085 Res = getDeviceAttrRaw(CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT, TmpInt);1086 if (Res == CUDA_SUCCESS)1087 Info.add("Number of Multiprocessors", TmpInt, "",1088 DeviceInfo::NUM_COMPUTE_UNITS);1089 1090 Res = getDeviceAttrRaw(CU_DEVICE_ATTRIBUTE_GPU_OVERLAP, TmpInt);1091 if (Res == CUDA_SUCCESS)1092 Info.add("Concurrent Copy and Execution", (bool)TmpInt);1093 1094 Res = getDeviceAttrRaw(CU_DEVICE_ATTRIBUTE_TOTAL_CONSTANT_MEMORY, TmpInt);1095 if (Res == CUDA_SUCCESS)1096 Info.add("Total Constant Memory", TmpInt, "bytes");1097 1098 Info.add("Max Shared Memory per Block", MaxBlockSharedMemSize, "bytes",1099 DeviceInfo::WORK_GROUP_LOCAL_MEM_SIZE);1100 1101 Res = getDeviceAttrRaw(CU_DEVICE_ATTRIBUTE_MAX_REGISTERS_PER_BLOCK, TmpInt);1102 if (Res == CUDA_SUCCESS)1103 Info.add("Registers per Block", TmpInt);1104 1105 Res = getDeviceAttrRaw(CU_DEVICE_ATTRIBUTE_WARP_SIZE, TmpInt);1106 if (Res == CUDA_SUCCESS)1107 Info.add("Warp Size", TmpInt);1108 1109 Res = getDeviceAttrRaw(CU_DEVICE_ATTRIBUTE_MAX_THREADS_PER_BLOCK, TmpInt);1110 if (Res == CUDA_SUCCESS)1111 Info.add("Maximum Threads per Block", TmpInt, "",1112 DeviceInfo::MAX_WORK_GROUP_SIZE);1113 1114 auto &MaxBlock = *Info.add("Maximum Block Dimensions", std::monostate{}, "",1115 DeviceInfo::MAX_WORK_GROUP_SIZE_PER_DIMENSION);1116 Res = getDeviceAttrRaw(CU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_X, TmpInt);1117 if (Res == CUDA_SUCCESS)1118 MaxBlock.add("x", TmpInt);1119 Res = getDeviceAttrRaw(CU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_Y, TmpInt);1120 if (Res == CUDA_SUCCESS)1121 MaxBlock.add("y", TmpInt);1122 Res = getDeviceAttrRaw(CU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_Z, TmpInt);1123 if (Res == CUDA_SUCCESS)1124 MaxBlock.add("z", TmpInt);1125 1126 // TODO: I assume CUDA devices have no limit on the amount of threads,1127 // verify this1128 Info.add("Maximum Grid Size", std::numeric_limits<uint32_t>::max(), "",1129 DeviceInfo::MAX_WORK_SIZE);1130 1131 auto &MaxGrid = *Info.add("Maximum Grid Dimensions", std::monostate{}, "",1132 DeviceInfo::MAX_WORK_SIZE_PER_DIMENSION);1133 Res = getDeviceAttrRaw(CU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_X, TmpInt);1134 if (Res == CUDA_SUCCESS)1135 MaxGrid.add("x", TmpInt);1136 Res = getDeviceAttrRaw(CU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_Y, TmpInt);1137 if (Res == CUDA_SUCCESS)1138 MaxGrid.add("y", TmpInt);1139 Res = getDeviceAttrRaw(CU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_Z, TmpInt);1140 if (Res == CUDA_SUCCESS)1141 MaxGrid.add("z", TmpInt);1142 1143 Res = getDeviceAttrRaw(CU_DEVICE_ATTRIBUTE_MAX_PITCH, TmpInt);1144 if (Res == CUDA_SUCCESS)1145 Info.add("Maximum Memory Pitch", TmpInt, "bytes");1146 1147 Res = getDeviceAttrRaw(CU_DEVICE_ATTRIBUTE_TEXTURE_ALIGNMENT, TmpInt);1148 if (Res == CUDA_SUCCESS)1149 Info.add("Texture Alignment", TmpInt, "bytes");1150 1151 Res = getDeviceAttrRaw(CU_DEVICE_ATTRIBUTE_CLOCK_RATE, TmpInt);1152 if (Res == CUDA_SUCCESS)1153 Info.add("Clock Rate", TmpInt / 1000, "MHz",1154 DeviceInfo::MAX_CLOCK_FREQUENCY);1155 1156 Res = getDeviceAttrRaw(CU_DEVICE_ATTRIBUTE_KERNEL_EXEC_TIMEOUT, TmpInt);1157 if (Res == CUDA_SUCCESS)1158 Info.add("Execution Timeout", (bool)TmpInt);1159 1160 Res = getDeviceAttrRaw(CU_DEVICE_ATTRIBUTE_INTEGRATED, TmpInt);1161 if (Res == CUDA_SUCCESS)1162 Info.add("Integrated Device", (bool)TmpInt);1163 1164 Res = getDeviceAttrRaw(CU_DEVICE_ATTRIBUTE_CAN_MAP_HOST_MEMORY, TmpInt);1165 if (Res == CUDA_SUCCESS)1166 Info.add("Can Map Host Memory", (bool)TmpInt);1167 1168 Res = getDeviceAttrRaw(CU_DEVICE_ATTRIBUTE_COMPUTE_MODE, TmpInt);1169 if (Res == CUDA_SUCCESS) {1170 if (TmpInt == CU_COMPUTEMODE_DEFAULT)1171 TmpCharPtr = "Default";1172 else if (TmpInt == CU_COMPUTEMODE_PROHIBITED)1173 TmpCharPtr = "Prohibited";1174 else if (TmpInt == CU_COMPUTEMODE_EXCLUSIVE_PROCESS)1175 TmpCharPtr = "Exclusive process";1176 else1177 TmpCharPtr = "Unknown";1178 Info.add("Compute Mode", TmpCharPtr);1179 }1180 1181 Res = getDeviceAttrRaw(CU_DEVICE_ATTRIBUTE_CONCURRENT_KERNELS, TmpInt);1182 if (Res == CUDA_SUCCESS)1183 Info.add("Concurrent Kernels", (bool)TmpInt);1184 1185 Res = getDeviceAttrRaw(CU_DEVICE_ATTRIBUTE_ECC_ENABLED, TmpInt);1186 if (Res == CUDA_SUCCESS)1187 Info.add("ECC Enabled", (bool)TmpInt);1188 1189 Res = getDeviceAttrRaw(CU_DEVICE_ATTRIBUTE_MEMORY_CLOCK_RATE, TmpInt);1190 if (Res == CUDA_SUCCESS)1191 Info.add("Memory Clock Rate", TmpInt / 1000, "MHz",1192 DeviceInfo::MEMORY_CLOCK_RATE);1193 1194 Res = getDeviceAttrRaw(CU_DEVICE_ATTRIBUTE_GLOBAL_MEMORY_BUS_WIDTH, TmpInt);1195 if (Res == CUDA_SUCCESS)1196 Info.add("Memory Bus Width", TmpInt, "bits");1197 1198 Res = getDeviceAttrRaw(CU_DEVICE_ATTRIBUTE_L2_CACHE_SIZE, TmpInt);1199 if (Res == CUDA_SUCCESS)1200 Info.add("L2 Cache Size", TmpInt, "bytes");1201 1202 Res = getDeviceAttrRaw(CU_DEVICE_ATTRIBUTE_MAX_THREADS_PER_MULTIPROCESSOR,1203 TmpInt);1204 if (Res == CUDA_SUCCESS)1205 Info.add("Max Threads Per SMP", TmpInt);1206 1207 Res = getDeviceAttrRaw(CU_DEVICE_ATTRIBUTE_ASYNC_ENGINE_COUNT, TmpInt);1208 if (Res == CUDA_SUCCESS)1209 Info.add("Async Engines", TmpInt);1210 1211 Res = getDeviceAttrRaw(CU_DEVICE_ATTRIBUTE_UNIFIED_ADDRESSING, TmpInt);1212 if (Res == CUDA_SUCCESS)1213 Info.add("Unified Addressing", (bool)TmpInt);1214 1215 Res = getDeviceAttrRaw(CU_DEVICE_ATTRIBUTE_MANAGED_MEMORY, TmpInt);1216 if (Res == CUDA_SUCCESS)1217 Info.add("Managed Memory", (bool)TmpInt);1218 1219 Res =1220 getDeviceAttrRaw(CU_DEVICE_ATTRIBUTE_CONCURRENT_MANAGED_ACCESS, TmpInt);1221 if (Res == CUDA_SUCCESS)1222 Info.add("Concurrent Managed Memory", (bool)TmpInt);1223 1224 Res = getDeviceAttrRaw(CU_DEVICE_ATTRIBUTE_COMPUTE_PREEMPTION_SUPPORTED,1225 TmpInt);1226 if (Res == CUDA_SUCCESS)1227 Info.add("Preemption Supported", (bool)TmpInt);1228 1229 Res = getDeviceAttrRaw(CU_DEVICE_ATTRIBUTE_COOPERATIVE_LAUNCH, TmpInt);1230 if (Res == CUDA_SUCCESS)1231 Info.add("Cooperative Launch", (bool)TmpInt);1232 1233 Res = getDeviceAttrRaw(CU_DEVICE_ATTRIBUTE_MULTI_GPU_BOARD, TmpInt);1234 if (Res == CUDA_SUCCESS)1235 Info.add("Multi-Device Boars", (bool)TmpInt);1236 1237 Info.add("Compute Capabilities", ComputeCapability.str());1238 1239 return Info;1240 }1241 1242 /// Getters and setters for stack and heap sizes.1243 Error getDeviceStackSize(uint64_t &Value) override {1244 return getCtxLimit(CU_LIMIT_STACK_SIZE, Value);1245 }1246 Error setDeviceStackSize(uint64_t Value) override {1247 return setCtxLimit(CU_LIMIT_STACK_SIZE, Value);1248 }1249 bool hasDeviceHeapSize() override { return true; }1250 Error getDeviceHeapSize(uint64_t &Value) override {1251 return getCtxLimit(CU_LIMIT_MALLOC_HEAP_SIZE, Value);1252 }1253 Error setDeviceHeapSize(uint64_t Value) override {1254 return setCtxLimit(CU_LIMIT_MALLOC_HEAP_SIZE, Value);1255 }1256 Error getDeviceMemorySize(uint64_t &Value) override {1257 CUresult Res = cuDeviceTotalMem(&Value, Device);1258 return Plugin::check(Res, "error in getDeviceMemorySize %s");1259 }1260 1261 /// CUDA-specific functions for getting and setting context limits.1262 Error setCtxLimit(CUlimit Kind, uint64_t Value) {1263 CUresult Res = cuCtxSetLimit(Kind, Value);1264 return Plugin::check(Res, "error in cuCtxSetLimit: %s");1265 }1266 Error getCtxLimit(CUlimit Kind, uint64_t &Value) {1267 CUresult Res = cuCtxGetLimit(&Value, Kind);1268 return Plugin::check(Res, "error in cuCtxGetLimit: %s");1269 }1270 1271 /// CUDA-specific function to get device attributes.1272 Error getDeviceAttr(uint32_t Kind, uint32_t &Value) {1273 // TODO: Warn if the new value is larger than the old.1274 CUresult Res =1275 cuDeviceGetAttribute((int *)&Value, (CUdevice_attribute)Kind, Device);1276 return Plugin::check(Res, "error in cuDeviceGetAttribute: %s");1277 }1278 1279 CUresult getDeviceAttrRaw(uint32_t Kind, int &Value) {1280 return cuDeviceGetAttribute(&Value, (CUdevice_attribute)Kind, Device);1281 }1282 1283 /// See GenericDeviceTy::getComputeUnitKind().1284 std::string getComputeUnitKind() const override {1285 return ComputeCapability.str();1286 }1287 1288 /// Returns the clock frequency for the given NVPTX device.1289 uint64_t getClockFrequency() const override { return 1000000000; }1290 1291private:1292 using CUDAStreamManagerTy = GenericDeviceResourceManagerTy<CUDAStreamRef>;1293 using CUDAEventManagerTy = GenericDeviceResourceManagerTy<CUDAEventRef>;1294 1295 Error callGlobalCtorDtorCommon(GenericPluginTy &Plugin, DeviceImageTy &Image,1296 bool IsCtor) {1297 const char *KernelName = IsCtor ? "nvptx$device$init" : "nvptx$device$fini";1298 // Perform a quick check for the named kernel in the image. The kernel1299 // should be created by the 'nvptx-lower-ctor-dtor' pass.1300 GenericGlobalHandlerTy &Handler = Plugin.getGlobalHandler();1301 if (!Handler.isSymbolInImage(*this, Image, KernelName))1302 return Plugin::success();1303 1304 // The Nvidia backend cannot handle creating the ctor / dtor array1305 // automatically so we must create it ourselves. The backend will emit1306 // several globals that contain function pointers we can call. These are1307 // prefixed with a known name due to Nvidia's lack of section support.1308 auto ELFObjOrErr = Handler.getELFObjectFile(Image);1309 if (!ELFObjOrErr)1310 return ELFObjOrErr.takeError();1311 1312 // Search for all symbols that contain a constructor or destructor.1313 SmallVector<std::pair<StringRef, uint16_t>> Funcs;1314 for (ELFSymbolRef Sym : (*ELFObjOrErr)->symbols()) {1315 auto NameOrErr = Sym.getName();1316 if (!NameOrErr)1317 return NameOrErr.takeError();1318 1319 if (!NameOrErr->starts_with(IsCtor ? "__init_array_object_"1320 : "__fini_array_object_"))1321 continue;1322 1323 uint16_t Priority;1324 if (NameOrErr->rsplit('_').second.getAsInteger(10, Priority))1325 return Plugin::error(ErrorCode::INVALID_BINARY,1326 "invalid priority for constructor or destructor");1327 1328 Funcs.emplace_back(*NameOrErr, Priority);1329 }1330 1331 // Sort the created array to be in priority order.1332 llvm::sort(Funcs, [=](auto X, auto Y) { return X.second < Y.second; });1333 1334 // Allocate a buffer to store all of the known constructor / destructor1335 // functions in so we can iterate them on the device.1336 auto BufferOrErr =1337 allocate(Funcs.size() * sizeof(void *), nullptr, TARGET_ALLOC_DEVICE);1338 if (!BufferOrErr)1339 return BufferOrErr.takeError();1340 1341 void *Buffer = *BufferOrErr;1342 if (!Buffer)1343 return Plugin::error(ErrorCode::OUT_OF_RESOURCES,1344 "failed to allocate memory for global buffer");1345 1346 auto *GlobalPtrStart = reinterpret_cast<uintptr_t *>(Buffer);1347 auto *GlobalPtrStop = reinterpret_cast<uintptr_t *>(Buffer) + Funcs.size();1348 1349 SmallVector<void *> FunctionPtrs(Funcs.size());1350 std::size_t Idx = 0;1351 for (auto [Name, Priority] : Funcs) {1352 GlobalTy FunctionAddr(Name.str(), sizeof(void *), &FunctionPtrs[Idx++]);1353 if (auto Err = Handler.readGlobalFromDevice(*this, Image, FunctionAddr))1354 return Err;1355 }1356 1357 // Copy the local buffer to the device.1358 if (auto Err = dataSubmit(GlobalPtrStart, FunctionPtrs.data(),1359 FunctionPtrs.size() * sizeof(void *), nullptr))1360 return Err;1361 1362 // Copy the created buffer to the appropriate symbols so the kernel can1363 // iterate through them.1364 GlobalTy StartGlobal(IsCtor ? "__init_array_start" : "__fini_array_start",1365 sizeof(void *), &GlobalPtrStart);1366 if (auto Err = Handler.writeGlobalToDevice(*this, Image, StartGlobal))1367 return Err;1368 1369 GlobalTy StopGlobal(IsCtor ? "__init_array_end" : "__fini_array_end",1370 sizeof(void *), &GlobalPtrStop);1371 if (auto Err = Handler.writeGlobalToDevice(*this, Image, StopGlobal))1372 return Err;1373 1374 CUDAKernelTy CUDAKernel(KernelName);1375 1376 if (auto Err = CUDAKernel.init(*this, Image))1377 return Err;1378 1379 AsyncInfoWrapperTy AsyncInfoWrapper(*this, nullptr);1380 1381 KernelArgsTy KernelArgs = {};1382 uint32_t NumBlocksAndThreads[3] = {1u, 1u, 1u};1383 if (auto Err = CUDAKernel.launchImpl(1384 *this, NumBlocksAndThreads, NumBlocksAndThreads, KernelArgs,1385 KernelLaunchParamsTy{}, AsyncInfoWrapper))1386 return Err;1387 1388 Error Err = Plugin::success();1389 AsyncInfoWrapper.finalize(Err);1390 if (Err)1391 return Err;1392 1393 return free(Buffer, TARGET_ALLOC_DEVICE);1394 }1395 1396 /// Stream manager for CUDA streams.1397 CUDAStreamManagerTy CUDAStreamManager;1398 1399 /// Event manager for CUDA events.1400 CUDAEventManagerTy CUDAEventManager;1401 1402 /// The device's context. This context should be set before performing1403 /// operations on the device.1404 CUcontext Context = nullptr;1405 1406 /// The CUDA device handler.1407 CUdevice Device = CU_DEVICE_INVALID;1408 1409 /// The memory mapped addresses and their handles1410 std::unordered_map<CUdeviceptr, CUmemGenericAllocationHandle> DeviceMMaps;1411 1412 /// The compute capability of the corresponding CUDA device.1413 struct ComputeCapabilityTy {1414 uint32_t Major;1415 uint32_t Minor;1416 std::string str() const {1417 return "sm_" + std::to_string(Major * 10 + Minor);1418 }1419 } ComputeCapability;1420 1421 /// The maximum number of warps that can be resident on all the SMs1422 /// simultaneously.1423 uint32_t HardwareParallelism = 0;1424};1425 1426Error CUDAKernelTy::launchImpl(GenericDeviceTy &GenericDevice,1427 uint32_t NumThreads[3], uint32_t NumBlocks[3],1428 KernelArgsTy &KernelArgs,1429 KernelLaunchParamsTy LaunchParams,1430 AsyncInfoWrapperTy &AsyncInfoWrapper) const {1431 CUDADeviceTy &CUDADevice = static_cast<CUDADeviceTy &>(GenericDevice);1432 1433 CUstream Stream;1434 if (auto Err = CUDADevice.getStream(AsyncInfoWrapper, Stream))1435 return Err;1436 1437 uint32_t MaxDynCGroupMem =1438 std::max(KernelArgs.DynCGroupMem, GenericDevice.getDynamicMemorySize());1439 1440 void *Config[] = {CU_LAUNCH_PARAM_BUFFER_POINTER, LaunchParams.Data,1441 CU_LAUNCH_PARAM_BUFFER_SIZE,1442 reinterpret_cast<void *>(&LaunchParams.Size),1443 CU_LAUNCH_PARAM_END};1444 1445 // If we are running an RPC server we want to wake up the server thread1446 // whenever there is a kernel running and let it sleep otherwise.1447 if (GenericDevice.getRPCServer())1448 GenericDevice.Plugin.getRPCServer().Thread->notify();1449 1450 // In case we require more memory than the current limit.1451 if (MaxDynCGroupMem >= MaxDynCGroupMemLimit) {1452 CUresult AttrResult = cuFuncSetAttribute(1453 Func, CU_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES, MaxDynCGroupMem);1454 if (auto Err = Plugin::check(1455 AttrResult,1456 "error in cuFuncSetAttribute while setting the memory limits: %s"))1457 return Err;1458 MaxDynCGroupMemLimit = MaxDynCGroupMem;1459 }1460 1461 CUresult Res = cuLaunchKernel(Func, NumBlocks[0], NumBlocks[1], NumBlocks[2],1462 NumThreads[0], NumThreads[1], NumThreads[2],1463 MaxDynCGroupMem, Stream, nullptr, Config);1464 1465 // Register a callback to indicate when the kernel is complete.1466 if (GenericDevice.getRPCServer())1467 cuLaunchHostFunc(1468 Stream,1469 [](void *Data) {1470 GenericPluginTy &Plugin = *reinterpret_cast<GenericPluginTy *>(Data);1471 Plugin.getRPCServer().Thread->finish();1472 },1473 &GenericDevice.Plugin);1474 1475 return Plugin::check(Res, "error in cuLaunchKernel for '%s': %s", getName());1476}1477 1478/// Class implementing the CUDA-specific functionalities of the global handler.1479class CUDAGlobalHandlerTy final : public GenericGlobalHandlerTy {1480public:1481 /// Get the metadata of a global from the device. The name and size of the1482 /// global is read from DeviceGlobal and the address of the global is written1483 /// to DeviceGlobal.1484 Error getGlobalMetadataFromDevice(GenericDeviceTy &Device,1485 DeviceImageTy &Image,1486 GlobalTy &DeviceGlobal) override {1487 CUDADeviceImageTy &CUDAImage = static_cast<CUDADeviceImageTy &>(Image);1488 1489 const char *GlobalName = DeviceGlobal.getName().data();1490 1491 size_t CUSize;1492 CUdeviceptr CUPtr;1493 CUresult Res =1494 cuModuleGetGlobal(&CUPtr, &CUSize, CUDAImage.getModule(), GlobalName);1495 if (auto Err = Plugin::check(Res, "error in cuModuleGetGlobal for '%s': %s",1496 GlobalName))1497 return Err;1498 1499 if (DeviceGlobal.getSize() && CUSize != DeviceGlobal.getSize())1500 return Plugin::error(1501 ErrorCode::INVALID_BINARY,1502 "failed to load global '%s' due to size mismatch (%zu != %zu)",1503 GlobalName, CUSize, (size_t)DeviceGlobal.getSize());1504 1505 DeviceGlobal.setPtr(reinterpret_cast<void *>(CUPtr));1506 DeviceGlobal.setSize(CUSize);1507 1508 return Plugin::success();1509 }1510};1511 1512/// Class implementing the CUDA-specific functionalities of the plugin.1513struct CUDAPluginTy final : public GenericPluginTy {1514 /// Create a CUDA plugin.1515 CUDAPluginTy() : GenericPluginTy(getTripleArch()) {}1516 1517 /// This class should not be copied.1518 CUDAPluginTy(const CUDAPluginTy &) = delete;1519 CUDAPluginTy(CUDAPluginTy &&) = delete;1520 1521 /// Initialize the plugin and return the number of devices.1522 Expected<int32_t> initImpl() override {1523 CUresult Res = cuInit(0);1524 if (Res == CUDA_ERROR_INVALID_HANDLE) {1525 // Cannot call cuGetErrorString if dlsym failed.1526 DP("Failed to load CUDA shared library\n");1527 return 0;1528 }1529 1530 if (Res == CUDA_ERROR_NO_DEVICE) {1531 // Do not initialize if there are no devices.1532 DP("There are no devices supporting CUDA.\n");1533 return 0;1534 }1535 1536 if (auto Err = Plugin::check(Res, "error in cuInit: %s"))1537 return std::move(Err);1538 1539 // Get the number of devices.1540 int NumDevices;1541 Res = cuDeviceGetCount(&NumDevices);1542 if (auto Err = Plugin::check(Res, "error in cuDeviceGetCount: %s"))1543 return std::move(Err);1544 1545 // Do not initialize if there are no devices.1546 if (NumDevices == 0)1547 DP("There are no devices supporting CUDA.\n");1548 1549 return NumDevices;1550 }1551 1552 /// Deinitialize the plugin.1553 Error deinitImpl() override { return Plugin::success(); }1554 1555 /// Creates a CUDA device to use for offloading.1556 GenericDeviceTy *createDevice(GenericPluginTy &Plugin, int32_t DeviceId,1557 int32_t NumDevices) override {1558 return new CUDADeviceTy(Plugin, DeviceId, NumDevices);1559 }1560 1561 /// Creates a CUDA global handler.1562 GenericGlobalHandlerTy *createGlobalHandler() override {1563 return new CUDAGlobalHandlerTy();1564 }1565 1566 /// Get the ELF code for recognizing the compatible image binary.1567 uint16_t getMagicElfBits() const override { return ELF::EM_CUDA; }1568 1569 Triple::ArchType getTripleArch() const override {1570 // TODO: I think we can drop the support for 32-bit NVPTX devices.1571 return Triple::nvptx64;1572 }1573 1574 const char *getName() const override { return GETNAME(TARGET_NAME); }1575 1576 /// Check whether the image is compatible with a CUDA device.1577 Expected<bool> isELFCompatible(uint32_t DeviceId,1578 StringRef Image) const override {1579 auto ElfOrErr =1580 ELF64LEObjectFile::create(MemoryBufferRef(Image, /*Identifier=*/""),1581 /*InitContent=*/false);1582 if (!ElfOrErr)1583 return ElfOrErr.takeError();1584 1585 // Get the numeric value for the image's `sm_` value.1586 const auto Header = ElfOrErr->getELFFile().getHeader();1587 unsigned SM =1588 Header.e_ident[ELF::EI_ABIVERSION] == ELF::ELFABIVERSION_CUDA_V11589 ? Header.e_flags & ELF::EF_CUDA_SM1590 : (Header.e_flags & ELF::EF_CUDA_SM_MASK) >> ELF::EF_CUDA_SM_OFFSET;1591 1592 CUdevice Device;1593 CUresult Res = cuDeviceGet(&Device, DeviceId);1594 if (auto Err = Plugin::check(Res, "error in cuDeviceGet: %s"))1595 return std::move(Err);1596 1597 int32_t Major, Minor;1598 Res = cuDeviceGetAttribute(1599 &Major, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, Device);1600 if (auto Err = Plugin::check(Res, "error in cuDeviceGetAttribute: %s"))1601 return std::move(Err);1602 1603 Res = cuDeviceGetAttribute(1604 &Minor, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, Device);1605 if (auto Err = Plugin::check(Res, "error in cuDeviceGetAttribute: %s"))1606 return std::move(Err);1607 1608 int32_t ImageMajor = SM / 10;1609 int32_t ImageMinor = SM % 10;1610 1611 // A cubin generated for a certain compute capability is supported to1612 // run on any GPU with the same major revision and same or higher minor1613 // revision.1614 return Major == ImageMajor && Minor >= ImageMinor;1615 }1616};1617 1618Error CUDADeviceTy::dataExchangeImpl(const void *SrcPtr,1619 GenericDeviceTy &DstGenericDevice,1620 void *DstPtr, int64_t Size,1621 AsyncInfoWrapperTy &AsyncInfoWrapper) {1622 if (auto Err = setContext())1623 return Err;1624 1625 CUDADeviceTy &DstDevice = static_cast<CUDADeviceTy &>(DstGenericDevice);1626 1627 CUresult Res;1628 int32_t DstDeviceId = DstDevice.DeviceId;1629 CUdeviceptr CUSrcPtr = (CUdeviceptr)SrcPtr;1630 CUdeviceptr CUDstPtr = (CUdeviceptr)DstPtr;1631 1632 int CanAccessPeer = 0;1633 if (DeviceId != DstDeviceId) {1634 // Make sure the lock is released before performing the copies.1635 std::lock_guard<std::mutex> Lock(PeerAccessesLock);1636 1637 switch (PeerAccesses[DstDeviceId]) {1638 case PeerAccessState::AVAILABLE:1639 CanAccessPeer = 1;1640 break;1641 case PeerAccessState::UNAVAILABLE:1642 CanAccessPeer = 0;1643 break;1644 case PeerAccessState::PENDING:1645 // Check whether the source device can access the destination device.1646 Res = cuDeviceCanAccessPeer(&CanAccessPeer, Device, DstDevice.Device);1647 if (auto Err = Plugin::check(Res, "Error in cuDeviceCanAccessPeer: %s"))1648 return Err;1649 1650 if (CanAccessPeer) {1651 Res = cuCtxEnablePeerAccess(DstDevice.Context, 0);1652 if (Res == CUDA_ERROR_TOO_MANY_PEERS) {1653 // Resources may be exhausted due to many P2P links.1654 CanAccessPeer = 0;1655 DP("Too many P2P so fall back to D2D memcpy");1656 } else if (auto Err =1657 Plugin::check(Res, "error in cuCtxEnablePeerAccess: %s"))1658 return Err;1659 }1660 PeerAccesses[DstDeviceId] = (CanAccessPeer)1661 ? PeerAccessState::AVAILABLE1662 : PeerAccessState::UNAVAILABLE;1663 }1664 }1665 1666 CUstream Stream;1667 if (auto Err = getStream(AsyncInfoWrapper, Stream))1668 return Err;1669 1670 if (CanAccessPeer) {1671 // TODO: Should we fallback to D2D if peer access fails?1672 Res = cuMemcpyPeerAsync(CUDstPtr, Context, CUSrcPtr, DstDevice.Context,1673 Size, Stream);1674 return Plugin::check(Res, "error in cuMemcpyPeerAsync: %s");1675 }1676 1677 // Fallback to D2D copy.1678 Res = cuMemcpyDtoDAsync(CUDstPtr, CUSrcPtr, Size, Stream);1679 return Plugin::check(Res, "error in cuMemcpyDtoDAsync: %s");1680}1681 1682template <typename... ArgsTy>1683static Error Plugin::check(int32_t Code, const char *ErrFmt, ArgsTy... Args) {1684 CUresult ResultCode = static_cast<CUresult>(Code);1685 if (ResultCode == CUDA_SUCCESS)1686 return Plugin::success();1687 1688 const char *Desc = "Unknown error";1689 CUresult Ret = cuGetErrorString(ResultCode, &Desc);1690 if (Ret != CUDA_SUCCESS)1691 REPORT("Unrecognized " GETNAME(TARGET_NAME) " error code %d\n", Code);1692 1693 // TODO: Add more entries to this switch1694 ErrorCode OffloadErrCode;1695 switch (ResultCode) {1696 case CUDA_ERROR_NOT_FOUND:1697 OffloadErrCode = ErrorCode::NOT_FOUND;1698 break;1699 default:1700 OffloadErrCode = ErrorCode::UNKNOWN;1701 }1702 1703 // TODO: Create a map for CUDA error codes to Offload error codes1704 return Plugin::error(OffloadErrCode, ErrFmt, Args..., Desc);1705}1706 1707} // namespace plugin1708} // namespace target1709} // namespace omp1710} // namespace llvm1711 1712extern "C" {1713llvm::omp::target::plugin::GenericPluginTy *createPlugin_cuda() {1714 return new llvm::omp::target::plugin::CUDAPluginTy();1715}1716}1717