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