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