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