clang/lib/CodeGen/CGCUDANV.cpp

   1 //===----- CGCUDANV.cpp - Interface to NVIDIA CUDA Runtime ----------------===//
   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-exception
   6 //
   7 //===----------------------------------------------------------------------===//
   8 //
   9 // This provides a class for CUDA code generation targeting the NVIDIA CUDA
  10 // runtime library.
  11 //
  12 //===----------------------------------------------------------------------===//
  13
  14 #include "CGCUDARuntime.h"
  15 #include "CGCXXABI.h"
  16 #include "CodeGenFunction.h"
  17 #include "CodeGenModule.h"
  18 #include "clang/AST/Decl.h"
  19 #include "clang/Basic/Cuda.h"
  20 #include "clang/CodeGen/CodeGenABITypes.h"
  21 #include "clang/CodeGen/ConstantInitBuilder.h"
  22 #include "llvm/Frontend/Offloading/Utility.h"
  23 #include "llvm/IR/BasicBlock.h"
  24 #include "llvm/IR/Constants.h"
  25 #include "llvm/IR/DerivedTypes.h"
  26 #include "llvm/IR/ReplaceConstant.h"
  27 #include "llvm/Support/Format.h"
  28 #include "llvm/Support/VirtualFileSystem.h"
  29
  30 using namespace clang;
  31 using namespace CodeGen;
  32
  33 namespace {
  34 constexpr unsigned CudaFatMagic = 0x466243b1;
  35 constexpr unsigned HIPFatMagic = 0x48495046; // "HIPF"
  36
  37 class CGNVCUDARuntime : public CGCUDARuntime {
  38
  39 private:
  40   llvm::IntegerType *IntTy, *SizeTy;
  41   llvm::Type *VoidTy;
  42   llvm::PointerType *PtrTy;
  43
  44   /// Convenience reference to LLVM Context
  45   llvm::LLVMContext &Context;
  46   /// Convenience reference to the current module
  47   llvm::Module &TheModule;
  48   /// Keeps track of kernel launch stubs and handles emitted in this module
  49   struct KernelInfo {
  50     llvm::Function *Kernel; // stub function to help launch kernel
  51     const Decl *D;
  52   };
  53   llvm::SmallVector<KernelInfo, 16> EmittedKernels;
  54   // Map a kernel mangled name to a symbol for identifying kernel in host code
  55   // For CUDA, the symbol for identifying the kernel is the same as the device
  56   // stub function. For HIP, they are different.
  57   llvm::DenseMap<StringRef, llvm::GlobalValue *> KernelHandles;
  58   // Map a kernel handle to the kernel stub.
  59   llvm::DenseMap<llvm::GlobalValue *, llvm::Function *> KernelStubs;
  60   struct VarInfo {
  61     llvm::GlobalVariable *Var;
  62     const VarDecl *D;
  63     DeviceVarFlags Flags;
  64   };
  65   llvm::SmallVector<VarInfo, 16> DeviceVars;
  66   /// Keeps track of variable containing handle of GPU binary. Populated by
  67   /// ModuleCtorFunction() and used to create corresponding cleanup calls in
  68   /// ModuleDtorFunction()
  69   llvm::GlobalVariable *GpuBinaryHandle = nullptr;
  70   /// Whether we generate relocatable device code.
  71   bool RelocatableDeviceCode;
  72   /// Mangle context for device.
  73   std::unique_ptr<MangleContext> DeviceMC;
  74   /// Some zeros used for GEPs.
  75   llvm::Constant *Zeros[2];
  76
  77   llvm::FunctionCallee getSetupArgumentFn() const;
  78   llvm::FunctionCallee getLaunchFn() const;
  79
  80   llvm::FunctionType *getRegisterGlobalsFnTy() const;
  81   llvm::FunctionType *getCallbackFnTy() const;
  82   llvm::FunctionType *getRegisterLinkedBinaryFnTy() const;
  83   std::string addPrefixToName(StringRef FuncName) const;
  84   std::string addUnderscoredPrefixToName(StringRef FuncName) const;
  85
  86   /// Creates a function to register all kernel stubs generated in this module.
  87   llvm::Function *makeRegisterGlobalsFn();
  88
  89   /// Helper function that generates a constant string and returns a pointer to
  90   /// the start of the string.  The result of this function can be used anywhere
  91   /// where the C code specifies const char*.
  92   llvm::Constant *makeConstantString(const std::string &Str,
  93                                      const std::string &Name = "") {
  94     auto ConstStr = CGM.GetAddrOfConstantCString(Str, Name.c_str());
  95     return llvm::ConstantExpr::getGetElementPtr(ConstStr.getElementType(),
  96                                                 ConstStr.getPointer(), Zeros);
  97   }
  98
  99   /// Helper function which generates an initialized constant array from Str,
 100   /// and optionally sets section name and alignment. AddNull specifies whether
 101   /// the array should nave NUL termination.
 102   llvm::Constant *makeConstantArray(StringRef Str,
 103                                     StringRef Name = "",
 104                                     StringRef SectionName = "",
 105                                     unsigned Alignment = 0,
 106                                     bool AddNull = false) {
 107     llvm::Constant *Value =
 108         llvm::ConstantDataArray::getString(Context, Str, AddNull);
 109     auto *GV = new llvm::GlobalVariable(
 110         TheModule, Value->getType(), /*isConstant=*/true,
 111         llvm::GlobalValue::PrivateLinkage, Value, Name);
 112     if (!SectionName.empty()) {
 113       GV->setSection(SectionName);
 114       // Mark the address as used which make sure that this section isn't
 115       // merged and we will really have it in the object file.
 116       GV->setUnnamedAddr(llvm::GlobalValue::UnnamedAddr::None);
 117     }
 118     if (Alignment)
 119       GV->setAlignment(llvm::Align(Alignment));
 120     return llvm::ConstantExpr::getGetElementPtr(GV->getValueType(), GV, Zeros);
 121   }
 122
 123   /// Helper function that generates an empty dummy function returning void.
 124   llvm::Function *makeDummyFunction(llvm::FunctionType *FnTy) {
 125     assert(FnTy->getReturnType()->isVoidTy() &&
 126            "Can only generate dummy functions returning void!");
 127     llvm::Function *DummyFunc = llvm::Function::Create(
 128         FnTy, llvm::GlobalValue::InternalLinkage, "dummy", &TheModule);
 129
 130     llvm::BasicBlock *DummyBlock =
 131         llvm::BasicBlock::Create(Context, "", DummyFunc);
 132     CGBuilderTy FuncBuilder(CGM, Context);
 133     FuncBuilder.SetInsertPoint(DummyBlock);
 134     FuncBuilder.CreateRetVoid();
 135
 136     return DummyFunc;
 137   }
 138
 139   void emitDeviceStubBodyLegacy(CodeGenFunction &CGF, FunctionArgList &Args);
 140   void emitDeviceStubBodyNew(CodeGenFunction &CGF, FunctionArgList &Args);
 141   std::string getDeviceSideName(const NamedDecl *ND) override;
 142
 143   void registerDeviceVar(const VarDecl *VD, llvm::GlobalVariable &Var,
 144                          bool Extern, bool Constant) {
 145     DeviceVars.push_back({&Var,
 146                           VD,
 147                           {DeviceVarFlags::Variable, Extern, Constant,
 148                            VD->hasAttr<HIPManagedAttr>(),
 149                            /*Normalized*/ false, 0}});
 150   }
 151   void registerDeviceSurf(const VarDecl *VD, llvm::GlobalVariable &Var,
 152                           bool Extern, int Type) {
 153     DeviceVars.push_back({&Var,
 154                           VD,
 155                           {DeviceVarFlags::Surface, Extern, /*Constant*/ false,
 156                            /*Managed*/ false,
 157                            /*Normalized*/ false, Type}});
 158   }
 159   void registerDeviceTex(const VarDecl *VD, llvm::GlobalVariable &Var,
 160                          bool Extern, int Type, bool Normalized) {
 161     DeviceVars.push_back({&Var,
 162                           VD,
 163                           {DeviceVarFlags::Texture, Extern, /*Constant*/ false,
 164                            /*Managed*/ false, Normalized, Type}});
 165   }
 166
 167   /// Creates module constructor function
 168   llvm::Function *makeModuleCtorFunction();
 169   /// Creates module destructor function
 170   llvm::Function *makeModuleDtorFunction();
 171   /// Transform managed variables for device compilation.
 172   void transformManagedVars();
 173   /// Create offloading entries to register globals in RDC mode.
 174   void createOffloadingEntries();
 175
 176 public:
 177   CGNVCUDARuntime(CodeGenModule &CGM);
 178
 179   llvm::GlobalValue *getKernelHandle(llvm::Function *F, GlobalDecl GD) override;
 180   llvm::Function *getKernelStub(llvm::GlobalValue *Handle) override {
 181     auto Loc = KernelStubs.find(Handle);
 182     assert(Loc != KernelStubs.end());
 183     return Loc->second;
 184   }
 185   void emitDeviceStub(CodeGenFunction &CGF, FunctionArgList &Args) override;
 186   void handleVarRegistration(const VarDecl *VD,
 187                              llvm::GlobalVariable &Var) override;
 188   void
 189   internalizeDeviceSideVar(const VarDecl *D,
 190                            llvm::GlobalValue::LinkageTypes &Linkage) override;
 191
 192   llvm::Function *finalizeModule() override;
 193 };
 194
 195 } // end anonymous namespace
 196
 197 std::string CGNVCUDARuntime::addPrefixToName(StringRef FuncName) const {
 198   if (CGM.getLangOpts().HIP)
 199     return ((Twine("hip") + Twine(FuncName)).str());
 200   return ((Twine("cuda") + Twine(FuncName)).str());
 201 }
 202 std::string
 203 CGNVCUDARuntime::addUnderscoredPrefixToName(StringRef FuncName) const {
 204   if (CGM.getLangOpts().HIP)
 205     return ((Twine("__hip") + Twine(FuncName)).str());
 206   return ((Twine("__cuda") + Twine(FuncName)).str());
 207 }
 208
 209 static std::unique_ptr<MangleContext> InitDeviceMC(CodeGenModule &CGM) {
 210   // If the host and device have different C++ ABIs, mark it as the device
 211   // mangle context so that the mangling needs to retrieve the additional
 212   // device lambda mangling number instead of the regular host one.
 213   if (CGM.getContext().getAuxTargetInfo() &&
 214       CGM.getContext().getTargetInfo().getCXXABI().isMicrosoft() &&
 215       CGM.getContext().getAuxTargetInfo()->getCXXABI().isItaniumFamily()) {
 216     return std::unique_ptr<MangleContext>(
 217         CGM.getContext().createDeviceMangleContext(
 218             *CGM.getContext().getAuxTargetInfo()));
 219   }
 220
 221   return std::unique_ptr<MangleContext>(CGM.getContext().createMangleContext(
 222       CGM.getContext().getAuxTargetInfo()));
 223 }
 224
 225 CGNVCUDARuntime::CGNVCUDARuntime(CodeGenModule &CGM)
 226     : CGCUDARuntime(CGM), Context(CGM.getLLVMContext()),
 227       TheModule(CGM.getModule()),
 228       RelocatableDeviceCode(CGM.getLangOpts().GPURelocatableDeviceCode),
 229       DeviceMC(InitDeviceMC(CGM)) {
 230   IntTy = CGM.IntTy;
 231   SizeTy = CGM.SizeTy;
 232   VoidTy = CGM.VoidTy;
 233   Zeros[0] = llvm::ConstantInt::get(SizeTy, 0);
 234   Zeros[1] = Zeros[0];
 235   PtrTy = CGM.UnqualPtrTy;
 236 }
 237
 238 llvm::FunctionCallee CGNVCUDARuntime::getSetupArgumentFn() const {
 239   // cudaError_t cudaSetupArgument(void *, size_t, size_t)
 240   llvm::Type *Params[] = {PtrTy, SizeTy, SizeTy};
 241   return CGM.CreateRuntimeFunction(
 242       llvm::FunctionType::get(IntTy, Params, false),
 243       addPrefixToName("SetupArgument"));
 244 }
 245
 246 llvm::FunctionCallee CGNVCUDARuntime::getLaunchFn() const {
 247   if (CGM.getLangOpts().HIP) {
 248     // hipError_t hipLaunchByPtr(char *);
 249     return CGM.CreateRuntimeFunction(
 250         llvm::FunctionType::get(IntTy, PtrTy, false), "hipLaunchByPtr");
 251   }
 252   // cudaError_t cudaLaunch(char *);
 253   return CGM.CreateRuntimeFunction(llvm::FunctionType::get(IntTy, PtrTy, false),
 254                                    "cudaLaunch");
 255 }
 256
 257 llvm::FunctionType *CGNVCUDARuntime::getRegisterGlobalsFnTy() const {
 258   return llvm::FunctionType::get(VoidTy, PtrTy, false);
 259 }
 260
 261 llvm::FunctionType *CGNVCUDARuntime::getCallbackFnTy() const {
 262   return llvm::FunctionType::get(VoidTy, PtrTy, false);
 263 }
 264
 265 llvm::FunctionType *CGNVCUDARuntime::getRegisterLinkedBinaryFnTy() const {
 266   llvm::Type *Params[] = {llvm::PointerType::getUnqual(Context), PtrTy, PtrTy,
 267                           llvm::PointerType::getUnqual(Context)};
 268   return llvm::FunctionType::get(VoidTy, Params, false);
 269 }
 270
 271 std::string CGNVCUDARuntime::getDeviceSideName(const NamedDecl *ND) {
 272   GlobalDecl GD;
 273   // D could be either a kernel or a variable.
 274   if (auto *FD = dyn_cast<FunctionDecl>(ND))
 275     GD = GlobalDecl(FD, KernelReferenceKind::Kernel);
 276   else
 277     GD = GlobalDecl(ND);
 278   std::string DeviceSideName;
 279   MangleContext *MC;
 280   if (CGM.getLangOpts().CUDAIsDevice)
 281     MC = &CGM.getCXXABI().getMangleContext();
 282   else
 283     MC = DeviceMC.get();
 284   if (MC->shouldMangleDeclName(ND)) {
 285     SmallString<256> Buffer;
 286     llvm::raw_svector_ostream Out(Buffer);
 287     MC->mangleName(GD, Out);
 288     DeviceSideName = std::string(Out.str());
 289   } else
 290     DeviceSideName = std::string(ND->getIdentifier()->getName());
 291
 292   // Make unique name for device side static file-scope variable for HIP.
 293   if (CGM.getContext().shouldExternalize(ND) &&
 294       CGM.getLangOpts().GPURelocatableDeviceCode) {
 295     SmallString<256> Buffer;
 296     llvm::raw_svector_ostream Out(Buffer);
 297     Out << DeviceSideName;
 298     CGM.printPostfixForExternalizedDecl(Out, ND);
 299     DeviceSideName = std::string(Out.str());
 300   }
 301   return DeviceSideName;
 302 }
 303
 304 void CGNVCUDARuntime::emitDeviceStub(CodeGenFunction &CGF,
 305                                      FunctionArgList &Args) {
 306   EmittedKernels.push_back({CGF.CurFn, CGF.CurFuncDecl});
 307   if (auto *GV =
 308           dyn_cast<llvm::GlobalVariable>(KernelHandles[CGF.CurFn->getName()])) {
 309     GV->setLinkage(CGF.CurFn->getLinkage());
 310     GV->setInitializer(CGF.CurFn);
 311   }
 312   if (CudaFeatureEnabled(CGM.getTarget().getSDKVersion(),
 313                          CudaFeature::CUDA_USES_NEW_LAUNCH) ||
 314       (CGF.getLangOpts().HIP && CGF.getLangOpts().HIPUseNewLaunchAPI))
 315     emitDeviceStubBodyNew(CGF, Args);
 316   else
 317     emitDeviceStubBodyLegacy(CGF, Args);
 318 }
 319
 320 // CUDA 9.0+ uses new way to launch kernels. Parameters are packed in a local
 321 // array and kernels are launched using cudaLaunchKernel().
 322 void CGNVCUDARuntime::emitDeviceStubBodyNew(CodeGenFunction &CGF,
 323                                             FunctionArgList &Args) {
 324   // Build the shadow stack entry at the very start of the function.
 325
 326   // Calculate amount of space we will need for all arguments.  If we have no
 327   // args, allocate a single pointer so we still have a valid pointer to the
 328   // argument array that we can pass to runtime, even if it will be unused.
 329   Address KernelArgs = CGF.CreateTempAlloca(
 330       PtrTy, CharUnits::fromQuantity(16), "kernel_args",
 331       llvm::ConstantInt::get(SizeTy, std::max<size_t>(1, Args.size())));
 332   // Store pointers to the arguments in a locally allocated launch_args.
 333   for (unsigned i = 0; i < Args.size(); ++i) {
 334     llvm::Value* VarPtr = CGF.GetAddrOfLocalVar(Args[i]).getPointer();
 335     llvm::Value *VoidVarPtr = CGF.Builder.CreatePointerCast(VarPtr, PtrTy);
 336     CGF.Builder.CreateDefaultAlignedStore(
 337         VoidVarPtr,
 338         CGF.Builder.CreateConstGEP1_32(PtrTy, KernelArgs.getPointer(), i));
 339   }
 340
 341   llvm::BasicBlock *EndBlock = CGF.createBasicBlock("setup.end");
 342
 343   // Lookup cudaLaunchKernel/hipLaunchKernel function.
 344   // HIP kernel launching API name depends on -fgpu-default-stream option. For
 345   // the default value 'legacy', it is hipLaunchKernel. For 'per-thread',
 346   // it is hipLaunchKernel_spt.
 347   // cudaError_t cudaLaunchKernel(const void *func, dim3 gridDim, dim3 blockDim,
 348   //                              void **args, size_t sharedMem,
 349   //                              cudaStream_t stream);
 350   // hipError_t hipLaunchKernel[_spt](const void *func, dim3 gridDim,
 351   //                                  dim3 blockDim, void **args,
 352   //                                  size_t sharedMem, hipStream_t stream);
 353   TranslationUnitDecl *TUDecl = CGM.getContext().getTranslationUnitDecl();
 354   DeclContext *DC = TranslationUnitDecl::castToDeclContext(TUDecl);
 355   std::string KernelLaunchAPI = "LaunchKernel";
 356   if (CGF.getLangOpts().GPUDefaultStream ==
 357       LangOptions::GPUDefaultStreamKind::PerThread) {
 358     if (CGF.getLangOpts().HIP)
 359       KernelLaunchAPI = KernelLaunchAPI + "_spt";
 360     else if (CGF.getLangOpts().CUDA)
 361       KernelLaunchAPI = KernelLaunchAPI + "_ptsz";
 362   }
 363   auto LaunchKernelName = addPrefixToName(KernelLaunchAPI);
 364   IdentifierInfo &cudaLaunchKernelII =
 365       CGM.getContext().Idents.get(LaunchKernelName);
 366   FunctionDecl *cudaLaunchKernelFD = nullptr;
 367   for (auto *Result : DC->lookup(&cudaLaunchKernelII)) {
 368     if (FunctionDecl *FD = dyn_cast<FunctionDecl>(Result))
 369       cudaLaunchKernelFD = FD;
 370   }
 371
 372   if (cudaLaunchKernelFD == nullptr) {
 373     CGM.Error(CGF.CurFuncDecl->getLocation(),
 374               "Can't find declaration for " + LaunchKernelName);
 375     return;
 376   }
 377   // Create temporary dim3 grid_dim, block_dim.
 378   ParmVarDecl *GridDimParam = cudaLaunchKernelFD->getParamDecl(1);
 379   QualType Dim3Ty = GridDimParam->getType();
 380   Address GridDim =
 381       CGF.CreateMemTemp(Dim3Ty, CharUnits::fromQuantity(8), "grid_dim");
 382   Address BlockDim =
 383       CGF.CreateMemTemp(Dim3Ty, CharUnits::fromQuantity(8), "block_dim");
 384   Address ShmemSize =
 385       CGF.CreateTempAlloca(SizeTy, CGM.getSizeAlign(), "shmem_size");
 386   Address Stream = CGF.CreateTempAlloca(PtrTy, CGM.getPointerAlign(), "stream");
 387   llvm::FunctionCallee cudaPopConfigFn = CGM.CreateRuntimeFunction(
 388       llvm::FunctionType::get(IntTy,
 389                               {/*gridDim=*/GridDim.getType(),
 390                                /*blockDim=*/BlockDim.getType(),
 391                                /*ShmemSize=*/ShmemSize.getType(),
 392                                /*Stream=*/Stream.getType()},
 393                               /*isVarArg=*/false),
 394       addUnderscoredPrefixToName("PopCallConfiguration"));
 395
 396   CGF.EmitRuntimeCallOrInvoke(cudaPopConfigFn,
 397                               {GridDim.getPointer(), BlockDim.getPointer(),
 398                                ShmemSize.getPointer(), Stream.getPointer()});
 399
 400   // Emit the call to cudaLaunch
 401   llvm::Value *Kernel =
 402       CGF.Builder.CreatePointerCast(KernelHandles[CGF.CurFn->getName()], PtrTy);
 403   CallArgList LaunchKernelArgs;
 404   LaunchKernelArgs.add(RValue::get(Kernel),
 405                        cudaLaunchKernelFD->getParamDecl(0)->getType());
 406   LaunchKernelArgs.add(RValue::getAggregate(GridDim), Dim3Ty);
 407   LaunchKernelArgs.add(RValue::getAggregate(BlockDim), Dim3Ty);
 408   LaunchKernelArgs.add(RValue::get(KernelArgs.getPointer()),
 409                        cudaLaunchKernelFD->getParamDecl(3)->getType());
 410   LaunchKernelArgs.add(RValue::get(CGF.Builder.CreateLoad(ShmemSize)),
 411                        cudaLaunchKernelFD->getParamDecl(4)->getType());
 412   LaunchKernelArgs.add(RValue::get(CGF.Builder.CreateLoad(Stream)),
 413                        cudaLaunchKernelFD->getParamDecl(5)->getType());
 414
 415   QualType QT = cudaLaunchKernelFD->getType();
 416   QualType CQT = QT.getCanonicalType();
 417   llvm::Type *Ty = CGM.getTypes().ConvertType(CQT);
 418   llvm::FunctionType *FTy = cast<llvm::FunctionType>(Ty);
 419
 420   const CGFunctionInfo &FI =
 421       CGM.getTypes().arrangeFunctionDeclaration(cudaLaunchKernelFD);
 422   llvm::FunctionCallee cudaLaunchKernelFn =
 423       CGM.CreateRuntimeFunction(FTy, LaunchKernelName);
 424   CGF.EmitCall(FI, CGCallee::forDirect(cudaLaunchKernelFn), ReturnValueSlot(),
 425                LaunchKernelArgs);
 426   CGF.EmitBranch(EndBlock);
 427
 428   CGF.EmitBlock(EndBlock);
 429 }
 430
 431 void CGNVCUDARuntime::emitDeviceStubBodyLegacy(CodeGenFunction &CGF,
 432                                                FunctionArgList &Args) {
 433   // Emit a call to cudaSetupArgument for each arg in Args.
 434   llvm::FunctionCallee cudaSetupArgFn = getSetupArgumentFn();
 435   llvm::BasicBlock *EndBlock = CGF.createBasicBlock("setup.end");
 436   CharUnits Offset = CharUnits::Zero();
 437   for (const VarDecl *A : Args) {
 438     auto TInfo = CGM.getContext().getTypeInfoInChars(A->getType());
 439     Offset = Offset.alignTo(TInfo.Align);
 440     llvm::Value *Args[] = {
 441         CGF.Builder.CreatePointerCast(CGF.GetAddrOfLocalVar(A).getPointer(),
 442                                       PtrTy),
 443         llvm::ConstantInt::get(SizeTy, TInfo.Width.getQuantity()),
 444         llvm::ConstantInt::get(SizeTy, Offset.getQuantity()),
 445     };
 446     llvm::CallBase *CB = CGF.EmitRuntimeCallOrInvoke(cudaSetupArgFn, Args);
 447     llvm::Constant *Zero = llvm::ConstantInt::get(IntTy, 0);
 448     llvm::Value *CBZero = CGF.Builder.CreateICmpEQ(CB, Zero);
 449     llvm::BasicBlock *NextBlock = CGF.createBasicBlock("setup.next");
 450     CGF.Builder.CreateCondBr(CBZero, NextBlock, EndBlock);
 451     CGF.EmitBlock(NextBlock);
 452     Offset += TInfo.Width;
 453   }
 454
 455   // Emit the call to cudaLaunch
 456   llvm::FunctionCallee cudaLaunchFn = getLaunchFn();
 457   llvm::Value *Arg =
 458       CGF.Builder.CreatePointerCast(KernelHandles[CGF.CurFn->getName()], PtrTy);
 459   CGF.EmitRuntimeCallOrInvoke(cudaLaunchFn, Arg);
 460   CGF.EmitBranch(EndBlock);
 461
 462   CGF.EmitBlock(EndBlock);
 463 }
 464
 465 // Replace the original variable Var with the address loaded from variable
 466 // ManagedVar populated by HIP runtime.
 467 static void replaceManagedVar(llvm::GlobalVariable *Var,
 468                               llvm::GlobalVariable *ManagedVar) {
 469   SmallVector<SmallVector<llvm::User *, 8>, 8> WorkList;
 470   for (auto &&VarUse : Var->uses()) {
 471     WorkList.push_back({VarUse.getUser()});
 472   }
 473   while (!WorkList.empty()) {
 474     auto &&WorkItem = WorkList.pop_back_val();
 475     auto *U = WorkItem.back();
 476     if (isa<llvm::ConstantExpr>(U)) {
 477       for (auto &&UU : U->uses()) {
 478         WorkItem.push_back(UU.getUser());
 479         WorkList.push_back(WorkItem);
 480         WorkItem.pop_back();
 481       }
 482       continue;
 483     }
 484     if (auto *I = dyn_cast<llvm::Instruction>(U)) {
 485       llvm::Value *OldV = Var;
 486       llvm::Instruction *NewV =
 487           new llvm::LoadInst(Var->getType(), ManagedVar, "ld.managed", false,
 488                              llvm::Align(Var->getAlignment()), I);
 489       WorkItem.pop_back();
 490       // Replace constant expressions directly or indirectly using the managed
 491       // variable with instructions.
 492       for (auto &&Op : WorkItem) {
 493         auto *CE = cast<llvm::ConstantExpr>(Op);
 494         auto *NewInst = CE->getAsInstruction(I);
 495         NewInst->replaceUsesOfWith(OldV, NewV);
 496         OldV = CE;
 497         NewV = NewInst;
 498       }
 499       I->replaceUsesOfWith(OldV, NewV);
 500     } else {
 501       llvm_unreachable("Invalid use of managed variable");
 502     }
 503   }
 504 }
 505
 506 /// Creates a function that sets up state on the host side for CUDA objects that
 507 /// have a presence on both the host and device sides. Specifically, registers
 508 /// the host side of kernel functions and device global variables with the CUDA
 509 /// runtime.
 510 /// \code
 511 /// void __cuda_register_globals(void** GpuBinaryHandle) {
 512 ///    __cudaRegisterFunction(GpuBinaryHandle,Kernel0,...);
 513 ///    ...
 514 ///    __cudaRegisterFunction(GpuBinaryHandle,KernelM,...);
 515 ///    __cudaRegisterVar(GpuBinaryHandle, GlobalVar0, ...);
 516 ///    ...
 517 ///    __cudaRegisterVar(GpuBinaryHandle, GlobalVarN, ...);
 518 /// }
 519 /// \endcode
 520 llvm::Function *CGNVCUDARuntime::makeRegisterGlobalsFn() {
 521   // No need to register anything
 522   if (EmittedKernels.empty() && DeviceVars.empty())
 523     return nullptr;
 524
 525   llvm::Function *RegisterKernelsFunc = llvm::Function::Create(
 526       getRegisterGlobalsFnTy(), llvm::GlobalValue::InternalLinkage,
 527       addUnderscoredPrefixToName("_register_globals"), &TheModule);
 528   llvm::BasicBlock *EntryBB =
 529       llvm::BasicBlock::Create(Context, "entry", RegisterKernelsFunc);
 530   CGBuilderTy Builder(CGM, Context);
 531   Builder.SetInsertPoint(EntryBB);
 532
 533   // void __cudaRegisterFunction(void **, const char *, char *, const char *,
 534   //                             int, uint3*, uint3*, dim3*, dim3*, int*)
 535   llvm::Type *RegisterFuncParams[] = {
 536       PtrTy, PtrTy, PtrTy, PtrTy, IntTy,
 537       PtrTy, PtrTy, PtrTy, PtrTy, llvm::PointerType::getUnqual(Context)};
 538   llvm::FunctionCallee RegisterFunc = CGM.CreateRuntimeFunction(
 539       llvm::FunctionType::get(IntTy, RegisterFuncParams, false),
 540       addUnderscoredPrefixToName("RegisterFunction"));
 541
 542   // Extract GpuBinaryHandle passed as the first argument passed to
 543   // __cuda_register_globals() and generate __cudaRegisterFunction() call for
 544   // each emitted kernel.
 545   llvm::Argument &GpuBinaryHandlePtr = *RegisterKernelsFunc->arg_begin();
 546   for (auto &&I : EmittedKernels) {
 547     llvm::Constant *KernelName =
 548         makeConstantString(getDeviceSideName(cast<NamedDecl>(I.D)));
 549     llvm::Constant *NullPtr = llvm::ConstantPointerNull::get(PtrTy);
 550     llvm::Value *Args[] = {
 551         &GpuBinaryHandlePtr,
 552         KernelHandles[I.Kernel->getName()],
 553         KernelName,
 554         KernelName,
 555         llvm::ConstantInt::get(IntTy, -1),
 556         NullPtr,
 557         NullPtr,
 558         NullPtr,
 559         NullPtr,
 560         llvm::ConstantPointerNull::get(llvm::PointerType::getUnqual(Context))};
 561     Builder.CreateCall(RegisterFunc, Args);
 562   }
 563
 564   llvm::Type *VarSizeTy = IntTy;
 565   // For HIP or CUDA 9.0+, device variable size is type of `size_t`.
 566   if (CGM.getLangOpts().HIP ||
 567       ToCudaVersion(CGM.getTarget().getSDKVersion()) >= CudaVersion::CUDA_90)
 568     VarSizeTy = SizeTy;
 569
 570   // void __cudaRegisterVar(void **, char *, char *, const char *,
 571   //                        int, int, int, int)
 572   llvm::Type *RegisterVarParams[] = {PtrTy, PtrTy,     PtrTy, PtrTy,
 573                                      IntTy, VarSizeTy, IntTy, IntTy};
 574   llvm::FunctionCallee RegisterVar = CGM.CreateRuntimeFunction(
 575       llvm::FunctionType::get(VoidTy, RegisterVarParams, false),
 576       addUnderscoredPrefixToName("RegisterVar"));
 577   // void __hipRegisterManagedVar(void **, char *, char *, const char *,
 578   //                              size_t, unsigned)
 579   llvm::Type *RegisterManagedVarParams[] = {PtrTy, PtrTy,     PtrTy,
 580                                             PtrTy, VarSizeTy, IntTy};
 581   llvm::FunctionCallee RegisterManagedVar = CGM.CreateRuntimeFunction(
 582       llvm::FunctionType::get(VoidTy, RegisterManagedVarParams, false),
 583       addUnderscoredPrefixToName("RegisterManagedVar"));
 584   // void __cudaRegisterSurface(void **, const struct surfaceReference *,
 585   //                            const void **, const char *, int, int);
 586   llvm::FunctionCallee RegisterSurf = CGM.CreateRuntimeFunction(
 587       llvm::FunctionType::get(
 588           VoidTy, {PtrTy, PtrTy, PtrTy, PtrTy, IntTy, IntTy}, false),
 589       addUnderscoredPrefixToName("RegisterSurface"));
 590   // void __cudaRegisterTexture(void **, const struct textureReference *,
 591   //                            const void **, const char *, int, int, int)
 592   llvm::FunctionCallee RegisterTex = CGM.CreateRuntimeFunction(
 593       llvm::FunctionType::get(
 594           VoidTy, {PtrTy, PtrTy, PtrTy, PtrTy, IntTy, IntTy, IntTy}, false),
 595       addUnderscoredPrefixToName("RegisterTexture"));
 596   for (auto &&Info : DeviceVars) {
 597     llvm::GlobalVariable *Var = Info.Var;
 598     assert((!Var->isDeclaration() || Info.Flags.isManaged()) &&
 599            "External variables should not show up here, except HIP managed "
 600            "variables");
 601     llvm::Constant *VarName = makeConstantString(getDeviceSideName(Info.D));
 602     switch (Info.Flags.getKind()) {
 603     case DeviceVarFlags::Variable: {
 604       uint64_t VarSize =
 605           CGM.getDataLayout().getTypeAllocSize(Var->getValueType());
 606       if (Info.Flags.isManaged()) {
 607         auto *ManagedVar = new llvm::GlobalVariable(
 608             CGM.getModule(), Var->getType(),
 609             /*isConstant=*/false, Var->getLinkage(),
 610             /*Init=*/Var->isDeclaration()
 611                 ? nullptr
 612                 : llvm::ConstantPointerNull::get(Var->getType()),
 613             /*Name=*/"", /*InsertBefore=*/nullptr,
 614             llvm::GlobalVariable::NotThreadLocal);
 615         ManagedVar->setDSOLocal(Var->isDSOLocal());
 616         ManagedVar->setVisibility(Var->getVisibility());
 617         ManagedVar->setExternallyInitialized(true);
 618         ManagedVar->takeName(Var);
 619         Var->setName(Twine(ManagedVar->getName() + ".managed"));
 620         replaceManagedVar(Var, ManagedVar);
 621         llvm::Value *Args[] = {
 622             &GpuBinaryHandlePtr,
 623             ManagedVar,
 624             Var,
 625             VarName,
 626             llvm::ConstantInt::get(VarSizeTy, VarSize),
 627             llvm::ConstantInt::get(IntTy, Var->getAlignment())};
 628         if (!Var->isDeclaration())
 629           Builder.CreateCall(RegisterManagedVar, Args);
 630       } else {
 631         llvm::Value *Args[] = {
 632             &GpuBinaryHandlePtr,
 633             Var,
 634             VarName,
 635             VarName,
 636             llvm::ConstantInt::get(IntTy, Info.Flags.isExtern()),
 637             llvm::ConstantInt::get(VarSizeTy, VarSize),
 638             llvm::ConstantInt::get(IntTy, Info.Flags.isConstant()),
 639             llvm::ConstantInt::get(IntTy, 0)};
 640         Builder.CreateCall(RegisterVar, Args);
 641       }
 642       break;
 643     }
 644     case DeviceVarFlags::Surface:
 645       Builder.CreateCall(
 646           RegisterSurf,
 647           {&GpuBinaryHandlePtr, Var, VarName, VarName,
 648            llvm::ConstantInt::get(IntTy, Info.Flags.getSurfTexType()),
 649            llvm::ConstantInt::get(IntTy, Info.Flags.isExtern())});
 650       break;
 651     case DeviceVarFlags::Texture:
 652       Builder.CreateCall(
 653           RegisterTex,
 654           {&GpuBinaryHandlePtr, Var, VarName, VarName,
 655            llvm::ConstantInt::get(IntTy, Info.Flags.getSurfTexType()),
 656            llvm::ConstantInt::get(IntTy, Info.Flags.isNormalized()),
 657            llvm::ConstantInt::get(IntTy, Info.Flags.isExtern())});
 658       break;
 659     }
 660   }
 661
 662   Builder.CreateRetVoid();
 663   return RegisterKernelsFunc;
 664 }
 665
 666 /// Creates a global constructor function for the module:
 667 ///
 668 /// For CUDA:
 669 /// \code
 670 /// void __cuda_module_ctor() {
 671 ///     Handle = __cudaRegisterFatBinary(GpuBinaryBlob);
 672 ///     __cuda_register_globals(Handle);
 673 /// }
 674 /// \endcode
 675 ///
 676 /// For HIP:
 677 /// \code
 678 /// void __hip_module_ctor() {
 679 ///     if (__hip_gpubin_handle == 0) {
 680 ///         __hip_gpubin_handle  = __hipRegisterFatBinary(GpuBinaryBlob);
 681 ///         __hip_register_globals(__hip_gpubin_handle);
 682 ///     }
 683 /// }
 684 /// \endcode
 685 llvm::Function *CGNVCUDARuntime::makeModuleCtorFunction() {
 686   bool IsHIP = CGM.getLangOpts().HIP;
 687   bool IsCUDA = CGM.getLangOpts().CUDA;
 688   // No need to generate ctors/dtors if there is no GPU binary.
 689   StringRef CudaGpuBinaryFileName = CGM.getCodeGenOpts().CudaGpuBinaryFileName;
 690   if (CudaGpuBinaryFileName.empty() && !IsHIP)
 691     return nullptr;
 692   if ((IsHIP || (IsCUDA && !RelocatableDeviceCode)) && EmittedKernels.empty() &&
 693       DeviceVars.empty())
 694     return nullptr;
 695
 696   // void __{cuda|hip}_register_globals(void* handle);
 697   llvm::Function *RegisterGlobalsFunc = makeRegisterGlobalsFn();
 698   // We always need a function to pass in as callback. Create a dummy
 699   // implementation if we don't need to register anything.
 700   if (RelocatableDeviceCode && !RegisterGlobalsFunc)
 701     RegisterGlobalsFunc = makeDummyFunction(getRegisterGlobalsFnTy());
 702
 703   // void ** __{cuda|hip}RegisterFatBinary(void *);
 704   llvm::FunctionCallee RegisterFatbinFunc = CGM.CreateRuntimeFunction(
 705       llvm::FunctionType::get(PtrTy, PtrTy, false),
 706       addUnderscoredPrefixToName("RegisterFatBinary"));
 707   // struct { int magic, int version, void * gpu_binary, void * dont_care };
 708   llvm::StructType *FatbinWrapperTy =
 709       llvm::StructType::get(IntTy, IntTy, PtrTy, PtrTy);
 710
 711   // Register GPU binary with the CUDA runtime, store returned handle in a
 712   // global variable and save a reference in GpuBinaryHandle to be cleaned up
 713   // in destructor on exit. Then associate all known kernels with the GPU binary
 714   // handle so CUDA runtime can figure out what to call on the GPU side.
 715   std::unique_ptr<llvm::MemoryBuffer> CudaGpuBinary = nullptr;
 716   if (!CudaGpuBinaryFileName.empty()) {
 717     auto VFS = CGM.getFileSystem();
 718     auto CudaGpuBinaryOrErr =
 719         VFS->getBufferForFile(CudaGpuBinaryFileName, -1, false);
 720     if (std::error_code EC = CudaGpuBinaryOrErr.getError()) {
 721       CGM.getDiags().Report(diag::err_cannot_open_file)
 722           << CudaGpuBinaryFileName << EC.message();
 723       return nullptr;
 724     }
 725     CudaGpuBinary = std::move(CudaGpuBinaryOrErr.get());
 726   }
 727
 728   llvm::Function *ModuleCtorFunc = llvm::Function::Create(
 729       llvm::FunctionType::get(VoidTy, false),
 730       llvm::GlobalValue::InternalLinkage,
 731       addUnderscoredPrefixToName("_module_ctor"), &TheModule);
 732   llvm::BasicBlock *CtorEntryBB =
 733       llvm::BasicBlock::Create(Context, "entry", ModuleCtorFunc);
 734   CGBuilderTy CtorBuilder(CGM, Context);
 735
 736   CtorBuilder.SetInsertPoint(CtorEntryBB);
 737
 738   const char *FatbinConstantName;
 739   const char *FatbinSectionName;
 740   const char *ModuleIDSectionName;
 741   StringRef ModuleIDPrefix;
 742   llvm::Constant *FatBinStr;
 743   unsigned FatMagic;
 744   if (IsHIP) {
 745     FatbinConstantName = ".hip_fatbin";
 746     FatbinSectionName = ".hipFatBinSegment";
 747
 748     ModuleIDSectionName = "__hip_module_id";
 749     ModuleIDPrefix = "__hip_";
 750
 751     if (CudaGpuBinary) {
 752       // If fatbin is available from early finalization, create a string
 753       // literal containing the fat binary loaded from the given file.
 754       const unsigned HIPCodeObjectAlign = 4096;
 755       FatBinStr = makeConstantArray(std::string(CudaGpuBinary->getBuffer()), "",
 756                                     FatbinConstantName, HIPCodeObjectAlign);
 757     } else {
 758       // If fatbin is not available, create an external symbol
 759       // __hip_fatbin in section .hip_fatbin. The external symbol is supposed
 760       // to contain the fat binary but will be populated somewhere else,
 761       // e.g. by lld through link script.
 762       FatBinStr = new llvm::GlobalVariable(
 763         CGM.getModule(), CGM.Int8Ty,
 764         /*isConstant=*/true, llvm::GlobalValue::ExternalLinkage, nullptr,
 765         "__hip_fatbin", nullptr,
 766         llvm::GlobalVariable::NotThreadLocal);
 767       cast<llvm::GlobalVariable>(FatBinStr)->setSection(FatbinConstantName);
 768     }
 769
 770     FatMagic = HIPFatMagic;
 771   } else {
 772     if (RelocatableDeviceCode)
 773       FatbinConstantName = CGM.getTriple().isMacOSX()
 774                                ? "__NV_CUDA,__nv_relfatbin"
 775                                : "__nv_relfatbin";
 776     else
 777       FatbinConstantName =
 778           CGM.getTriple().isMacOSX() ? "__NV_CUDA,__nv_fatbin" : ".nv_fatbin";
 779     // NVIDIA's cuobjdump looks for fatbins in this section.
 780     FatbinSectionName =
 781         CGM.getTriple().isMacOSX() ? "__NV_CUDA,__fatbin" : ".nvFatBinSegment";
 782
 783     ModuleIDSectionName = CGM.getTriple().isMacOSX()
 784                               ? "__NV_CUDA,__nv_module_id"
 785                               : "__nv_module_id";
 786     ModuleIDPrefix = "__nv_";
 787
 788     // For CUDA, create a string literal containing the fat binary loaded from
 789     // the given file.
 790     FatBinStr = makeConstantArray(std::string(CudaGpuBinary->getBuffer()), "",
 791                                   FatbinConstantName, 8);
 792     FatMagic = CudaFatMagic;
 793   }
 794
 795   // Create initialized wrapper structure that points to the loaded GPU binary
 796   ConstantInitBuilder Builder(CGM);
 797   auto Values = Builder.beginStruct(FatbinWrapperTy);
 798   // Fatbin wrapper magic.
 799   Values.addInt(IntTy, FatMagic);
 800   // Fatbin version.
 801   Values.addInt(IntTy, 1);
 802   // Data.
 803   Values.add(FatBinStr);
 804   // Unused in fatbin v1.
 805   Values.add(llvm::ConstantPointerNull::get(PtrTy));
 806   llvm::GlobalVariable *FatbinWrapper = Values.finishAndCreateGlobal(
 807       addUnderscoredPrefixToName("_fatbin_wrapper"), CGM.getPointerAlign(),
 808       /*constant*/ true);
 809   FatbinWrapper->setSection(FatbinSectionName);
 810
 811   // There is only one HIP fat binary per linked module, however there are
 812   // multiple constructor functions. Make sure the fat binary is registered
 813   // only once. The constructor functions are executed by the dynamic loader
 814   // before the program gains control. The dynamic loader cannot execute the
 815   // constructor functions concurrently since doing that would not guarantee
 816   // thread safety of the loaded program. Therefore we can assume sequential
 817   // execution of constructor functions here.
 818   if (IsHIP) {
 819     auto Linkage = CudaGpuBinary ? llvm::GlobalValue::InternalLinkage :
 820         llvm::GlobalValue::LinkOnceAnyLinkage;
 821     llvm::BasicBlock *IfBlock =
 822         llvm::BasicBlock::Create(Context, "if", ModuleCtorFunc);
 823     llvm::BasicBlock *ExitBlock =
 824         llvm::BasicBlock::Create(Context, "exit", ModuleCtorFunc);
 825     // The name, size, and initialization pattern of this variable is part
 826     // of HIP ABI.
 827     GpuBinaryHandle = new llvm::GlobalVariable(
 828         TheModule, PtrTy, /*isConstant=*/false, Linkage,
 829         /*Initializer=*/llvm::ConstantPointerNull::get(PtrTy),
 830         "__hip_gpubin_handle");
 831     if (Linkage == llvm::GlobalValue::LinkOnceAnyLinkage)
 832       GpuBinaryHandle->setComdat(
 833           CGM.getModule().getOrInsertComdat(GpuBinaryHandle->getName()));
 834     GpuBinaryHandle->setAlignment(CGM.getPointerAlign().getAsAlign());
 835     // Prevent the weak symbol in different shared libraries being merged.
 836     if (Linkage != llvm::GlobalValue::InternalLinkage)
 837       GpuBinaryHandle->setVisibility(llvm::GlobalValue::HiddenVisibility);
 838     Address GpuBinaryAddr(
 839         GpuBinaryHandle, PtrTy,
 840         CharUnits::fromQuantity(GpuBinaryHandle->getAlignment()));
 841     {
 842       auto *HandleValue = CtorBuilder.CreateLoad(GpuBinaryAddr);
 843       llvm::Constant *Zero =
 844           llvm::Constant::getNullValue(HandleValue->getType());
 845       llvm::Value *EQZero = CtorBuilder.CreateICmpEQ(HandleValue, Zero);
 846       CtorBuilder.CreateCondBr(EQZero, IfBlock, ExitBlock);
 847     }
 848     {
 849       CtorBuilder.SetInsertPoint(IfBlock);
 850       // GpuBinaryHandle = __hipRegisterFatBinary(&FatbinWrapper);
 851       llvm::CallInst *RegisterFatbinCall =
 852           CtorBuilder.CreateCall(RegisterFatbinFunc, FatbinWrapper);
 853       CtorBuilder.CreateStore(RegisterFatbinCall, GpuBinaryAddr);
 854       CtorBuilder.CreateBr(ExitBlock);
 855     }
 856     {
 857       CtorBuilder.SetInsertPoint(ExitBlock);
 858       // Call __hip_register_globals(GpuBinaryHandle);
 859       if (RegisterGlobalsFunc) {
 860         auto *HandleValue = CtorBuilder.CreateLoad(GpuBinaryAddr);
 861         CtorBuilder.CreateCall(RegisterGlobalsFunc, HandleValue);
 862       }
 863     }
 864   } else if (!RelocatableDeviceCode) {
 865     // Register binary with CUDA runtime. This is substantially different in
 866     // default mode vs. separate compilation!
 867     // GpuBinaryHandle = __cudaRegisterFatBinary(&FatbinWrapper);
 868     llvm::CallInst *RegisterFatbinCall =
 869         CtorBuilder.CreateCall(RegisterFatbinFunc, FatbinWrapper);
 870     GpuBinaryHandle = new llvm::GlobalVariable(
 871         TheModule, PtrTy, false, llvm::GlobalValue::InternalLinkage,
 872         llvm::ConstantPointerNull::get(PtrTy), "__cuda_gpubin_handle");
 873     GpuBinaryHandle->setAlignment(CGM.getPointerAlign().getAsAlign());
 874     CtorBuilder.CreateAlignedStore(RegisterFatbinCall, GpuBinaryHandle,
 875                                    CGM.getPointerAlign());
 876
 877     // Call __cuda_register_globals(GpuBinaryHandle);
 878     if (RegisterGlobalsFunc)
 879       CtorBuilder.CreateCall(RegisterGlobalsFunc, RegisterFatbinCall);
 880
 881     // Call __cudaRegisterFatBinaryEnd(Handle) if this CUDA version needs it.
 882     if (CudaFeatureEnabled(CGM.getTarget().getSDKVersion(),
 883                            CudaFeature::CUDA_USES_FATBIN_REGISTER_END)) {
 884       // void __cudaRegisterFatBinaryEnd(void **);
 885       llvm::FunctionCallee RegisterFatbinEndFunc = CGM.CreateRuntimeFunction(
 886           llvm::FunctionType::get(VoidTy, PtrTy, false),
 887           "__cudaRegisterFatBinaryEnd");
 888       CtorBuilder.CreateCall(RegisterFatbinEndFunc, RegisterFatbinCall);
 889     }
 890   } else {
 891     // Generate a unique module ID.
 892     SmallString<64> ModuleID;
 893     llvm::raw_svector_ostream OS(ModuleID);
 894     OS << ModuleIDPrefix << llvm::format("%" PRIx64, FatbinWrapper->getGUID());
 895     llvm::Constant *ModuleIDConstant = makeConstantArray(
 896         std::string(ModuleID), "", ModuleIDSectionName, 32, /*AddNull=*/true);
 897
 898     // Create an alias for the FatbinWrapper that nvcc will look for.
 899     llvm::GlobalAlias::create(llvm::GlobalValue::ExternalLinkage,
 900                               Twine("__fatbinwrap") + ModuleID, FatbinWrapper);
 901
 902     // void __cudaRegisterLinkedBinary%ModuleID%(void (*)(void *), void *,
 903     // void *, void (*)(void **))
 904     SmallString<128> RegisterLinkedBinaryName("__cudaRegisterLinkedBinary");
 905     RegisterLinkedBinaryName += ModuleID;
 906     llvm::FunctionCallee RegisterLinkedBinaryFunc = CGM.CreateRuntimeFunction(
 907         getRegisterLinkedBinaryFnTy(), RegisterLinkedBinaryName);
 908
 909     assert(RegisterGlobalsFunc && "Expecting at least dummy function!");
 910     llvm::Value *Args[] = {RegisterGlobalsFunc, FatbinWrapper, ModuleIDConstant,
 911                            makeDummyFunction(getCallbackFnTy())};
 912     CtorBuilder.CreateCall(RegisterLinkedBinaryFunc, Args);
 913   }
 914
 915   // Create destructor and register it with atexit() the way NVCC does it. Doing
 916   // it during regular destructor phase worked in CUDA before 9.2 but results in
 917   // double-free in 9.2.
 918   if (llvm::Function *CleanupFn = makeModuleDtorFunction()) {
 919     // extern "C" int atexit(void (*f)(void));
 920     llvm::FunctionType *AtExitTy =
 921         llvm::FunctionType::get(IntTy, CleanupFn->getType(), false);
 922     llvm::FunctionCallee AtExitFunc =
 923         CGM.CreateRuntimeFunction(AtExitTy, "atexit", llvm::AttributeList(),
 924                                   /*Local=*/true);
 925     CtorBuilder.CreateCall(AtExitFunc, CleanupFn);
 926   }
 927
 928   CtorBuilder.CreateRetVoid();
 929   return ModuleCtorFunc;
 930 }
 931
 932 /// Creates a global destructor function that unregisters the GPU code blob
 933 /// registered by constructor.
 934 ///
 935 /// For CUDA:
 936 /// \code
 937 /// void __cuda_module_dtor() {
 938 ///     __cudaUnregisterFatBinary(Handle);
 939 /// }
 940 /// \endcode
 941 ///
 942 /// For HIP:
 943 /// \code
 944 /// void __hip_module_dtor() {
 945 ///     if (__hip_gpubin_handle) {
 946 ///         __hipUnregisterFatBinary(__hip_gpubin_handle);
 947 ///         __hip_gpubin_handle = 0;
 948 ///     }
 949 /// }
 950 /// \endcode
 951 llvm::Function *CGNVCUDARuntime::makeModuleDtorFunction() {
 952   // No need for destructor if we don't have a handle to unregister.
 953   if (!GpuBinaryHandle)
 954     return nullptr;
 955
 956   // void __cudaUnregisterFatBinary(void ** handle);
 957   llvm::FunctionCallee UnregisterFatbinFunc = CGM.CreateRuntimeFunction(
 958       llvm::FunctionType::get(VoidTy, PtrTy, false),
 959       addUnderscoredPrefixToName("UnregisterFatBinary"));
 960
 961   llvm::Function *ModuleDtorFunc = llvm::Function::Create(
 962       llvm::FunctionType::get(VoidTy, false),
 963       llvm::GlobalValue::InternalLinkage,
 964       addUnderscoredPrefixToName("_module_dtor"), &TheModule);
 965
 966   llvm::BasicBlock *DtorEntryBB =
 967       llvm::BasicBlock::Create(Context, "entry", ModuleDtorFunc);
 968   CGBuilderTy DtorBuilder(CGM, Context);
 969   DtorBuilder.SetInsertPoint(DtorEntryBB);
 970
 971   Address GpuBinaryAddr(
 972       GpuBinaryHandle, GpuBinaryHandle->getValueType(),
 973       CharUnits::fromQuantity(GpuBinaryHandle->getAlignment()));
 974   auto *HandleValue = DtorBuilder.CreateLoad(GpuBinaryAddr);
 975   // There is only one HIP fat binary per linked module, however there are
 976   // multiple destructor functions. Make sure the fat binary is unregistered
 977   // only once.
 978   if (CGM.getLangOpts().HIP) {
 979     llvm::BasicBlock *IfBlock =
 980         llvm::BasicBlock::Create(Context, "if", ModuleDtorFunc);
 981     llvm::BasicBlock *ExitBlock =
 982         llvm::BasicBlock::Create(Context, "exit", ModuleDtorFunc);
 983     llvm::Constant *Zero = llvm::Constant::getNullValue(HandleValue->getType());
 984     llvm::Value *NEZero = DtorBuilder.CreateICmpNE(HandleValue, Zero);
 985     DtorBuilder.CreateCondBr(NEZero, IfBlock, ExitBlock);
 986
 987     DtorBuilder.SetInsertPoint(IfBlock);
 988     DtorBuilder.CreateCall(UnregisterFatbinFunc, HandleValue);
 989     DtorBuilder.CreateStore(Zero, GpuBinaryAddr);
 990     DtorBuilder.CreateBr(ExitBlock);
 991
 992     DtorBuilder.SetInsertPoint(ExitBlock);
 993   } else {
 994     DtorBuilder.CreateCall(UnregisterFatbinFunc, HandleValue);
 995   }
 996   DtorBuilder.CreateRetVoid();
 997   return ModuleDtorFunc;
 998 }
 999
1000 CGCUDARuntime *CodeGen::CreateNVCUDARuntime(CodeGenModule &CGM) {
1001   return new CGNVCUDARuntime(CGM);
1002 }
1003
1004 void CGNVCUDARuntime::internalizeDeviceSideVar(
1005     const VarDecl *D, llvm::GlobalValue::LinkageTypes &Linkage) {
1006   // For -fno-gpu-rdc, host-side shadows of external declarations of device-side
1007   // global variables become internal definitions. These have to be internal in
1008   // order to prevent name conflicts with global host variables with the same
1009   // name in a different TUs.
1010   //
1011   // For -fgpu-rdc, the shadow variables should not be internalized because
1012   // they may be accessed by different TU.
1013   if (CGM.getLangOpts().GPURelocatableDeviceCode)
1014     return;
1015
1016   // __shared__ variables are odd. Shadows do get created, but
1017   // they are not registered with the CUDA runtime, so they
1018   // can't really be used to access their device-side
1019   // counterparts. It's not clear yet whether it's nvcc's bug or
1020   // a feature, but we've got to do the same for compatibility.
1021   if (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>() ||
1022       D->hasAttr<CUDASharedAttr>() ||
1023       D->getType()->isCUDADeviceBuiltinSurfaceType() ||
1024       D->getType()->isCUDADeviceBuiltinTextureType()) {
1025     Linkage = llvm::GlobalValue::InternalLinkage;
1026   }
1027 }
1028
1029 void CGNVCUDARuntime::handleVarRegistration(const VarDecl *D,
1030                                             llvm::GlobalVariable &GV) {
1031   if (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>()) {
1032     // Shadow variables and their properties must be registered with CUDA
1033     // runtime. Skip Extern global variables, which will be registered in
1034     // the TU where they are defined.
1035     //
1036     // Don't register a C++17 inline variable. The local symbol can be
1037     // discarded and referencing a discarded local symbol from outside the
1038     // comdat (__cuda_register_globals) is disallowed by the ELF spec.
1039     //
1040     // HIP managed variables need to be always recorded in device and host
1041     // compilations for transformation.
1042     //
1043     // HIP managed variables and variables in CUDADeviceVarODRUsedByHost are
1044     // added to llvm.compiler-used, therefore they are safe to be registered.
1045     if ((!D->hasExternalStorage() && !D->isInline()) ||
1046         CGM.getContext().CUDADeviceVarODRUsedByHost.contains(D) ||
1047         D->hasAttr<HIPManagedAttr>()) {
1048       registerDeviceVar(D, GV, !D->hasDefinition(),
1049                         D->hasAttr<CUDAConstantAttr>());
1050     }
1051   } else if (D->getType()->isCUDADeviceBuiltinSurfaceType() ||
1052              D->getType()->isCUDADeviceBuiltinTextureType()) {
1053     // Builtin surfaces and textures and their template arguments are
1054     // also registered with CUDA runtime.
1055     const auto *TD = cast<ClassTemplateSpecializationDecl>(
1056         D->getType()->castAs<RecordType>()->getDecl());
1057     const TemplateArgumentList &Args = TD->getTemplateArgs();
1058     if (TD->hasAttr<CUDADeviceBuiltinSurfaceTypeAttr>()) {
1059       assert(Args.size() == 2 &&
1060              "Unexpected number of template arguments of CUDA device "
1061              "builtin surface type.");
1062       auto SurfType = Args[1].getAsIntegral();
1063       if (!D->hasExternalStorage())
1064         registerDeviceSurf(D, GV, !D->hasDefinition(), SurfType.getSExtValue());
1065     } else {
1066       assert(Args.size() == 3 &&
1067              "Unexpected number of template arguments of CUDA device "
1068              "builtin texture type.");
1069       auto TexType = Args[1].getAsIntegral();
1070       auto Normalized = Args[2].getAsIntegral();
1071       if (!D->hasExternalStorage())
1072         registerDeviceTex(D, GV, !D->hasDefinition(), TexType.getSExtValue(),
1073                           Normalized.getZExtValue());
1074     }
1075   }
1076 }
1077
1078 // Transform managed variables to pointers to managed variables in device code.
1079 // Each use of the original managed variable is replaced by a load from the
1080 // transformed managed variable. The transformed managed variable contains
1081 // the address of managed memory which will be allocated by the runtime.
1082 void CGNVCUDARuntime::transformManagedVars() {
1083   for (auto &&Info : DeviceVars) {
1084     llvm::GlobalVariable *Var = Info.Var;
1085     if (Info.Flags.getKind() == DeviceVarFlags::Variable &&
1086         Info.Flags.isManaged()) {
1087       auto *ManagedVar = new llvm::GlobalVariable(
1088           CGM.getModule(), Var->getType(),
1089           /*isConstant=*/false, Var->getLinkage(),
1090           /*Init=*/Var->isDeclaration()
1091               ? nullptr
1092               : llvm::ConstantPointerNull::get(Var->getType()),
1093           /*Name=*/"", /*InsertBefore=*/nullptr,
1094           llvm::GlobalVariable::NotThreadLocal,
1095           CGM.getContext().getTargetAddressSpace(LangAS::cuda_device));
1096       ManagedVar->setDSOLocal(Var->isDSOLocal());
1097       ManagedVar->setVisibility(Var->getVisibility());
1098       ManagedVar->setExternallyInitialized(true);
1099       replaceManagedVar(Var, ManagedVar);
1100       ManagedVar->takeName(Var);
1101       Var->setName(Twine(ManagedVar->getName()) + ".managed");
1102       // Keep managed variables even if they are not used in device code since
1103       // they need to be allocated by the runtime.
1104       if (!Var->isDeclaration()) {
1105         assert(!ManagedVar->isDeclaration());
1106         CGM.addCompilerUsedGlobal(Var);
1107         CGM.addCompilerUsedGlobal(ManagedVar);
1108       }
1109     }
1110   }
1111 }
1112
1113 // Creates offloading entries for all the kernels and globals that must be
1114 // registered. The linker will provide a pointer to this section so we can
1115 // register the symbols with the linked device image.
1116 void CGNVCUDARuntime::createOffloadingEntries() {
1117   StringRef Section = CGM.getLangOpts().HIP ? "hip_offloading_entries"
1118                                             : "cuda_offloading_entries";
1119   llvm::Module &M = CGM.getModule();
1120   for (KernelInfo &I : EmittedKernels)
1121     llvm::offloading::emitOffloadingEntry(
1122         M, KernelHandles[I.Kernel->getName()],
1123         getDeviceSideName(cast<NamedDecl>(I.D)), /*Flags=*/0, /*Data=*/0,
1124         llvm::offloading::OffloadGlobalEntry, Section);
1125
1126   for (VarInfo &I : DeviceVars) {
1127     uint64_t VarSize =
1128         CGM.getDataLayout().getTypeAllocSize(I.Var->getValueType());
1129     int32_t Flags =
1130         (I.Flags.isExtern()
1131              ? static_cast<int32_t>(llvm::offloading::OffloadGlobalExtern)
1132              : 0) |
1133         (I.Flags.isConstant()
1134              ? static_cast<int32_t>(llvm::offloading::OffloadGlobalConstant)
1135              : 0) |
1136         (I.Flags.isNormalized()
1137              ? static_cast<int32_t>(llvm::offloading::OffloadGlobalNormalized)
1138              : 0);
1139     if (I.Flags.getKind() == DeviceVarFlags::Variable) {
1140       llvm::offloading::emitOffloadingEntry(
1141           M, I.Var, getDeviceSideName(I.D), VarSize,
1142           (I.Flags.isManaged() ? llvm::offloading::OffloadGlobalManagedEntry
1143                                : llvm::offloading::OffloadGlobalEntry) |
1144               Flags,
1145           /*Data=*/0, Section);
1146     } else if (I.Flags.getKind() == DeviceVarFlags::Surface) {
1147       llvm::offloading::emitOffloadingEntry(
1148           M, I.Var, getDeviceSideName(I.D), VarSize,
1149           llvm::offloading::OffloadGlobalSurfaceEntry | Flags,
1150           I.Flags.getSurfTexType(), Section);
1151     } else if (I.Flags.getKind() == DeviceVarFlags::Texture) {
1152       llvm::offloading::emitOffloadingEntry(
1153           M, I.Var, getDeviceSideName(I.D), VarSize,
1154           llvm::offloading::OffloadGlobalTextureEntry | Flags,
1155           I.Flags.getSurfTexType(), Section);
1156     }
1157   }
1158 }
1159
1160 // Returns module constructor to be added.
1161 llvm::Function *CGNVCUDARuntime::finalizeModule() {
1162   if (CGM.getLangOpts().CUDAIsDevice) {
1163     transformManagedVars();
1164
1165     // Mark ODR-used device variables as compiler used to prevent it from being
1166     // eliminated by optimization. This is necessary for device variables
1167     // ODR-used by host functions. Sema correctly marks them as ODR-used no
1168     // matter whether they are ODR-used by device or host functions.
1169     //
1170     // We do not need to do this if the variable has used attribute since it
1171     // has already been added.
1172     //
1173     // Static device variables have been externalized at this point, therefore
1174     // variables with LLVM private or internal linkage need not be added.
1175     for (auto &&Info : DeviceVars) {
1176       auto Kind = Info.Flags.getKind();
1177       if (!Info.Var->isDeclaration() &&
1178           !llvm::GlobalValue::isLocalLinkage(Info.Var->getLinkage()) &&
1179           (Kind == DeviceVarFlags::Variable ||
1180            Kind == DeviceVarFlags::Surface ||
1181            Kind == DeviceVarFlags::Texture) &&
1182           Info.D->isUsed() && !Info.D->hasAttr<UsedAttr>()) {
1183         CGM.addCompilerUsedGlobal(Info.Var);
1184       }
1185     }
1186     return nullptr;
1187   }
1188   if (CGM.getLangOpts().OffloadingNewDriver && RelocatableDeviceCode)
1189     createOffloadingEntries();
1190   else
1191     return makeModuleCtorFunction();
1192
1193   return nullptr;
1194 }
1195
1196 llvm::GlobalValue *CGNVCUDARuntime::getKernelHandle(llvm::Function *F,
1197                                                     GlobalDecl GD) {
1198   auto Loc = KernelHandles.find(F->getName());
1199   if (Loc != KernelHandles.end()) {
1200     auto OldHandle = Loc->second;
1201     if (KernelStubs[OldHandle] == F)
1202       return OldHandle;
1203
1204     // We've found the function name, but F itself has changed, so we need to
1205     // update the references.
1206     if (CGM.getLangOpts().HIP) {
1207       // For HIP compilation the handle itself does not change, so we only need
1208       // to update the Stub value.
1209       KernelStubs[OldHandle] = F;
1210       return OldHandle;
1211     }
1212     // For non-HIP compilation, erase the old Stub and fall-through to creating
1213     // new entries.
1214     KernelStubs.erase(OldHandle);
1215   }
1216
1217   if (!CGM.getLangOpts().HIP) {
1218     KernelHandles[F->getName()] = F;
1219     KernelStubs[F] = F;
1220     return F;
1221   }
1222
1223   auto *Var = new llvm::GlobalVariable(
1224       TheModule, F->getType(), /*isConstant=*/true, F->getLinkage(),
1225       /*Initializer=*/nullptr,
1226       CGM.getMangledName(
1227           GD.getWithKernelReferenceKind(KernelReferenceKind::Kernel)));
1228   Var->setAlignment(CGM.getPointerAlign().getAsAlign());
1229   Var->setDSOLocal(F->isDSOLocal());
1230   Var->setVisibility(F->getVisibility());
1231   auto *FD = cast<FunctionDecl>(GD.getDecl());
1232   auto *FT = FD->getPrimaryTemplate();
1233   if (!FT || FT->isThisDeclarationADefinition())
1234     CGM.maybeSetTrivialComdat(*FD, *Var);
1235   KernelHandles[F->getName()] = Var;
1236   KernelStubs[Var] = F;
1237   return Var;
1238 }