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 "CodeGenFunction.h"
16 #include "CodeGenModule.h"
17 #include "clang/AST/Decl.h"
18 #include "clang/Basic/Cuda.h"
19 #include "clang/CodeGen/CodeGenABITypes.h"
20 #include "clang/CodeGen/ConstantInitBuilder.h"
21 #include "llvm/IR/BasicBlock.h"
22 #include "llvm/IR/Constants.h"
23 #include "llvm/IR/DerivedTypes.h"
24 #include "llvm/Support/Format.h"
25 
26 using namespace clang;
27 using namespace CodeGen;
28 
29 namespace {
30 constexpr unsigned CudaFatMagic = 0x466243b1;
31 constexpr unsigned HIPFatMagic = 0x48495046; // "HIPF"
32 
33 class CGNVCUDARuntime : public CGCUDARuntime {
34 
35 private:
36   llvm::IntegerType *IntTy, *SizeTy;
37   llvm::Type *VoidTy;
38   llvm::PointerType *CharPtrTy, *VoidPtrTy, *VoidPtrPtrTy;
39 
40   /// Convenience reference to LLVM Context
41   llvm::LLVMContext &Context;
42   /// Convenience reference to the current module
43   llvm::Module &TheModule;
44   /// Keeps track of kernel launch stubs emitted in this module
45   llvm::SmallVector<llvm::Function *, 16> EmittedKernels;
46   llvm::SmallVector<std::pair<llvm::GlobalVariable *, unsigned>, 16> DeviceVars;
47   /// Keeps track of variable containing handle of GPU binary. Populated by
48   /// ModuleCtorFunction() and used to create corresponding cleanup calls in
49   /// ModuleDtorFunction()
50   llvm::GlobalVariable *GpuBinaryHandle = nullptr;
51   /// Whether we generate relocatable device code.
52   bool RelocatableDeviceCode;
53 
54   llvm::FunctionCallee getSetupArgumentFn() const;
55   llvm::FunctionCallee getLaunchFn() const;
56 
57   llvm::FunctionType *getRegisterGlobalsFnTy() const;
58   llvm::FunctionType *getCallbackFnTy() const;
59   llvm::FunctionType *getRegisterLinkedBinaryFnTy() const;
60   std::string addPrefixToName(StringRef FuncName) const;
61   std::string addUnderscoredPrefixToName(StringRef FuncName) const;
62 
63   /// Creates a function to register all kernel stubs generated in this module.
64   llvm::Function *makeRegisterGlobalsFn();
65 
66   /// Helper function that generates a constant string and returns a pointer to
67   /// the start of the string.  The result of this function can be used anywhere
68   /// where the C code specifies const char*.
69   llvm::Constant *makeConstantString(const std::string &Str,
70                                      const std::string &Name = "",
71                                      const std::string &SectionName = "",
72                                      unsigned Alignment = 0) {
73     llvm::Constant *Zeros[] = {llvm::ConstantInt::get(SizeTy, 0),
74                                llvm::ConstantInt::get(SizeTy, 0)};
75     auto ConstStr = CGM.GetAddrOfConstantCString(Str, Name.c_str());
76     llvm::GlobalVariable *GV =
77         cast<llvm::GlobalVariable>(ConstStr.getPointer());
78     if (!SectionName.empty()) {
79       GV->setSection(SectionName);
80       // Mark the address as used which make sure that this section isn't
81       // merged and we will really have it in the object file.
82       GV->setUnnamedAddr(llvm::GlobalValue::UnnamedAddr::None);
83     }
84     if (Alignment)
85       GV->setAlignment(Alignment);
86 
87     return llvm::ConstantExpr::getGetElementPtr(ConstStr.getElementType(),
88                                                 ConstStr.getPointer(), Zeros);
89   }
90 
91   /// Helper function that generates an empty dummy function returning void.
92   llvm::Function *makeDummyFunction(llvm::FunctionType *FnTy) {
93     assert(FnTy->getReturnType()->isVoidTy() &&
94            "Can only generate dummy functions returning void!");
95     llvm::Function *DummyFunc = llvm::Function::Create(
96         FnTy, llvm::GlobalValue::InternalLinkage, "dummy", &TheModule);
97 
98     llvm::BasicBlock *DummyBlock =
99         llvm::BasicBlock::Create(Context, "", DummyFunc);
100     CGBuilderTy FuncBuilder(CGM, Context);
101     FuncBuilder.SetInsertPoint(DummyBlock);
102     FuncBuilder.CreateRetVoid();
103 
104     return DummyFunc;
105   }
106 
107   void emitDeviceStubBodyLegacy(CodeGenFunction &CGF, FunctionArgList &Args);
108   void emitDeviceStubBodyNew(CodeGenFunction &CGF, FunctionArgList &Args);
109 
110 public:
111   CGNVCUDARuntime(CodeGenModule &CGM);
112 
113   void emitDeviceStub(CodeGenFunction &CGF, FunctionArgList &Args) override;
114   void registerDeviceVar(llvm::GlobalVariable &Var, unsigned Flags) override {
115     DeviceVars.push_back(std::make_pair(&Var, Flags));
116   }
117 
118   /// Creates module constructor function
119   llvm::Function *makeModuleCtorFunction() override;
120   /// Creates module destructor function
121   llvm::Function *makeModuleDtorFunction() override;
122 };
123 
124 }
125 
126 std::string CGNVCUDARuntime::addPrefixToName(StringRef FuncName) const {
127   if (CGM.getLangOpts().HIP)
128     return ((Twine("hip") + Twine(FuncName)).str());
129   return ((Twine("cuda") + Twine(FuncName)).str());
130 }
131 std::string
132 CGNVCUDARuntime::addUnderscoredPrefixToName(StringRef FuncName) const {
133   if (CGM.getLangOpts().HIP)
134     return ((Twine("__hip") + Twine(FuncName)).str());
135   return ((Twine("__cuda") + Twine(FuncName)).str());
136 }
137 
138 CGNVCUDARuntime::CGNVCUDARuntime(CodeGenModule &CGM)
139     : CGCUDARuntime(CGM), Context(CGM.getLLVMContext()),
140       TheModule(CGM.getModule()),
141       RelocatableDeviceCode(CGM.getLangOpts().GPURelocatableDeviceCode) {
142   CodeGen::CodeGenTypes &Types = CGM.getTypes();
143   ASTContext &Ctx = CGM.getContext();
144 
145   IntTy = CGM.IntTy;
146   SizeTy = CGM.SizeTy;
147   VoidTy = CGM.VoidTy;
148 
149   CharPtrTy = llvm::PointerType::getUnqual(Types.ConvertType(Ctx.CharTy));
150   VoidPtrTy = cast<llvm::PointerType>(Types.ConvertType(Ctx.VoidPtrTy));
151   VoidPtrPtrTy = VoidPtrTy->getPointerTo();
152 }
153 
154 llvm::FunctionCallee CGNVCUDARuntime::getSetupArgumentFn() const {
155   // cudaError_t cudaSetupArgument(void *, size_t, size_t)
156   llvm::Type *Params[] = {VoidPtrTy, SizeTy, SizeTy};
157   return CGM.CreateRuntimeFunction(
158       llvm::FunctionType::get(IntTy, Params, false),
159       addPrefixToName("SetupArgument"));
160 }
161 
162 llvm::FunctionCallee CGNVCUDARuntime::getLaunchFn() const {
163   if (CGM.getLangOpts().HIP) {
164     // hipError_t hipLaunchByPtr(char *);
165     return CGM.CreateRuntimeFunction(
166         llvm::FunctionType::get(IntTy, CharPtrTy, false), "hipLaunchByPtr");
167   } else {
168     // cudaError_t cudaLaunch(char *);
169     return CGM.CreateRuntimeFunction(
170         llvm::FunctionType::get(IntTy, CharPtrTy, false), "cudaLaunch");
171   }
172 }
173 
174 llvm::FunctionType *CGNVCUDARuntime::getRegisterGlobalsFnTy() const {
175   return llvm::FunctionType::get(VoidTy, VoidPtrPtrTy, false);
176 }
177 
178 llvm::FunctionType *CGNVCUDARuntime::getCallbackFnTy() const {
179   return llvm::FunctionType::get(VoidTy, VoidPtrTy, false);
180 }
181 
182 llvm::FunctionType *CGNVCUDARuntime::getRegisterLinkedBinaryFnTy() const {
183   auto CallbackFnTy = getCallbackFnTy();
184   auto RegisterGlobalsFnTy = getRegisterGlobalsFnTy();
185   llvm::Type *Params[] = {RegisterGlobalsFnTy->getPointerTo(), VoidPtrTy,
186                           VoidPtrTy, CallbackFnTy->getPointerTo()};
187   return llvm::FunctionType::get(VoidTy, Params, false);
188 }
189 
190 void CGNVCUDARuntime::emitDeviceStub(CodeGenFunction &CGF,
191                                      FunctionArgList &Args) {
192   EmittedKernels.push_back(CGF.CurFn);
193   if (CudaFeatureEnabled(CGM.getTarget().getSDKVersion(),
194                          CudaFeature::CUDA_USES_NEW_LAUNCH))
195     emitDeviceStubBodyNew(CGF, Args);
196   else
197     emitDeviceStubBodyLegacy(CGF, Args);
198 }
199 
200 // CUDA 9.0+ uses new way to launch kernels. Parameters are packed in a local
201 // array and kernels are launched using cudaLaunchKernel().
202 void CGNVCUDARuntime::emitDeviceStubBodyNew(CodeGenFunction &CGF,
203                                             FunctionArgList &Args) {
204   // Build the shadow stack entry at the very start of the function.
205 
206   // Calculate amount of space we will need for all arguments.  If we have no
207   // args, allocate a single pointer so we still have a valid pointer to the
208   // argument array that we can pass to runtime, even if it will be unused.
209   Address KernelArgs = CGF.CreateTempAlloca(
210       VoidPtrTy, CharUnits::fromQuantity(16), "kernel_args",
211       llvm::ConstantInt::get(SizeTy, std::max<size_t>(1, Args.size())));
212   // Store pointers to the arguments in a locally allocated launch_args.
213   for (unsigned i = 0; i < Args.size(); ++i) {
214     llvm::Value* VarPtr = CGF.GetAddrOfLocalVar(Args[i]).getPointer();
215     llvm::Value *VoidVarPtr = CGF.Builder.CreatePointerCast(VarPtr, VoidPtrTy);
216     CGF.Builder.CreateDefaultAlignedStore(
217         VoidVarPtr, CGF.Builder.CreateConstGEP1_32(KernelArgs.getPointer(), i));
218   }
219 
220   llvm::BasicBlock *EndBlock = CGF.createBasicBlock("setup.end");
221 
222   // Lookup cudaLaunchKernel function.
223   // cudaError_t cudaLaunchKernel(const void *func, dim3 gridDim, dim3 blockDim,
224   //                              void **args, size_t sharedMem,
225   //                              cudaStream_t stream);
226   TranslationUnitDecl *TUDecl = CGM.getContext().getTranslationUnitDecl();
227   DeclContext *DC = TranslationUnitDecl::castToDeclContext(TUDecl);
228   IdentifierInfo &cudaLaunchKernelII =
229       CGM.getContext().Idents.get("cudaLaunchKernel");
230   FunctionDecl *cudaLaunchKernelFD = nullptr;
231   for (const auto &Result : DC->lookup(&cudaLaunchKernelII)) {
232     if (FunctionDecl *FD = dyn_cast<FunctionDecl>(Result))
233       cudaLaunchKernelFD = FD;
234   }
235 
236   if (cudaLaunchKernelFD == nullptr) {
237     CGM.Error(CGF.CurFuncDecl->getLocation(),
238               "Can't find declaration for cudaLaunchKernel()");
239     return;
240   }
241   // Create temporary dim3 grid_dim, block_dim.
242   ParmVarDecl *GridDimParam = cudaLaunchKernelFD->getParamDecl(1);
243   QualType Dim3Ty = GridDimParam->getType();
244   Address GridDim =
245       CGF.CreateMemTemp(Dim3Ty, CharUnits::fromQuantity(8), "grid_dim");
246   Address BlockDim =
247       CGF.CreateMemTemp(Dim3Ty, CharUnits::fromQuantity(8), "block_dim");
248   Address ShmemSize =
249       CGF.CreateTempAlloca(SizeTy, CGM.getSizeAlign(), "shmem_size");
250   Address Stream =
251       CGF.CreateTempAlloca(VoidPtrTy, CGM.getPointerAlign(), "stream");
252   llvm::FunctionCallee cudaPopConfigFn = CGM.CreateRuntimeFunction(
253       llvm::FunctionType::get(IntTy,
254                               {/*gridDim=*/GridDim.getType(),
255                                /*blockDim=*/BlockDim.getType(),
256                                /*ShmemSize=*/ShmemSize.getType(),
257                                /*Stream=*/Stream.getType()},
258                               /*isVarArg=*/false),
259       "__cudaPopCallConfiguration");
260 
261   CGF.EmitRuntimeCallOrInvoke(cudaPopConfigFn,
262                               {GridDim.getPointer(), BlockDim.getPointer(),
263                                ShmemSize.getPointer(), Stream.getPointer()});
264 
265   // Emit the call to cudaLaunch
266   llvm::Value *Kernel = CGF.Builder.CreatePointerCast(CGF.CurFn, VoidPtrTy);
267   CallArgList LaunchKernelArgs;
268   LaunchKernelArgs.add(RValue::get(Kernel),
269                        cudaLaunchKernelFD->getParamDecl(0)->getType());
270   LaunchKernelArgs.add(RValue::getAggregate(GridDim), Dim3Ty);
271   LaunchKernelArgs.add(RValue::getAggregate(BlockDim), Dim3Ty);
272   LaunchKernelArgs.add(RValue::get(KernelArgs.getPointer()),
273                        cudaLaunchKernelFD->getParamDecl(3)->getType());
274   LaunchKernelArgs.add(RValue::get(CGF.Builder.CreateLoad(ShmemSize)),
275                        cudaLaunchKernelFD->getParamDecl(4)->getType());
276   LaunchKernelArgs.add(RValue::get(CGF.Builder.CreateLoad(Stream)),
277                        cudaLaunchKernelFD->getParamDecl(5)->getType());
278 
279   QualType QT = cudaLaunchKernelFD->getType();
280   QualType CQT = QT.getCanonicalType();
281   llvm::Type *Ty = CGM.getTypes().ConvertType(CQT);
282   llvm::FunctionType *FTy = dyn_cast<llvm::FunctionType>(Ty);
283 
284   const CGFunctionInfo &FI =
285       CGM.getTypes().arrangeFunctionDeclaration(cudaLaunchKernelFD);
286   llvm::FunctionCallee cudaLaunchKernelFn =
287       CGM.CreateRuntimeFunction(FTy, "cudaLaunchKernel");
288   CGF.EmitCall(FI, CGCallee::forDirect(cudaLaunchKernelFn), ReturnValueSlot(),
289                LaunchKernelArgs);
290   CGF.EmitBranch(EndBlock);
291 
292   CGF.EmitBlock(EndBlock);
293 }
294 
295 void CGNVCUDARuntime::emitDeviceStubBodyLegacy(CodeGenFunction &CGF,
296                                                FunctionArgList &Args) {
297   // Emit a call to cudaSetupArgument for each arg in Args.
298   llvm::FunctionCallee cudaSetupArgFn = getSetupArgumentFn();
299   llvm::BasicBlock *EndBlock = CGF.createBasicBlock("setup.end");
300   CharUnits Offset = CharUnits::Zero();
301   for (const VarDecl *A : Args) {
302     CharUnits TyWidth, TyAlign;
303     std::tie(TyWidth, TyAlign) =
304         CGM.getContext().getTypeInfoInChars(A->getType());
305     Offset = Offset.alignTo(TyAlign);
306     llvm::Value *Args[] = {
307         CGF.Builder.CreatePointerCast(CGF.GetAddrOfLocalVar(A).getPointer(),
308                                       VoidPtrTy),
309         llvm::ConstantInt::get(SizeTy, TyWidth.getQuantity()),
310         llvm::ConstantInt::get(SizeTy, Offset.getQuantity()),
311     };
312     llvm::CallBase *CB = CGF.EmitRuntimeCallOrInvoke(cudaSetupArgFn, Args);
313     llvm::Constant *Zero = llvm::ConstantInt::get(IntTy, 0);
314     llvm::Value *CBZero = CGF.Builder.CreateICmpEQ(CB, Zero);
315     llvm::BasicBlock *NextBlock = CGF.createBasicBlock("setup.next");
316     CGF.Builder.CreateCondBr(CBZero, NextBlock, EndBlock);
317     CGF.EmitBlock(NextBlock);
318     Offset += TyWidth;
319   }
320 
321   // Emit the call to cudaLaunch
322   llvm::FunctionCallee cudaLaunchFn = getLaunchFn();
323   llvm::Value *Arg = CGF.Builder.CreatePointerCast(CGF.CurFn, CharPtrTy);
324   CGF.EmitRuntimeCallOrInvoke(cudaLaunchFn, Arg);
325   CGF.EmitBranch(EndBlock);
326 
327   CGF.EmitBlock(EndBlock);
328 }
329 
330 /// Creates a function that sets up state on the host side for CUDA objects that
331 /// have a presence on both the host and device sides. Specifically, registers
332 /// the host side of kernel functions and device global variables with the CUDA
333 /// runtime.
334 /// \code
335 /// void __cuda_register_globals(void** GpuBinaryHandle) {
336 ///    __cudaRegisterFunction(GpuBinaryHandle,Kernel0,...);
337 ///    ...
338 ///    __cudaRegisterFunction(GpuBinaryHandle,KernelM,...);
339 ///    __cudaRegisterVar(GpuBinaryHandle, GlobalVar0, ...);
340 ///    ...
341 ///    __cudaRegisterVar(GpuBinaryHandle, GlobalVarN, ...);
342 /// }
343 /// \endcode
344 llvm::Function *CGNVCUDARuntime::makeRegisterGlobalsFn() {
345   // No need to register anything
346   if (EmittedKernels.empty() && DeviceVars.empty())
347     return nullptr;
348 
349   llvm::Function *RegisterKernelsFunc = llvm::Function::Create(
350       getRegisterGlobalsFnTy(), llvm::GlobalValue::InternalLinkage,
351       addUnderscoredPrefixToName("_register_globals"), &TheModule);
352   llvm::BasicBlock *EntryBB =
353       llvm::BasicBlock::Create(Context, "entry", RegisterKernelsFunc);
354   CGBuilderTy Builder(CGM, Context);
355   Builder.SetInsertPoint(EntryBB);
356 
357   // void __cudaRegisterFunction(void **, const char *, char *, const char *,
358   //                             int, uint3*, uint3*, dim3*, dim3*, int*)
359   llvm::Type *RegisterFuncParams[] = {
360       VoidPtrPtrTy, CharPtrTy, CharPtrTy, CharPtrTy, IntTy,
361       VoidPtrTy,    VoidPtrTy, VoidPtrTy, VoidPtrTy, IntTy->getPointerTo()};
362   llvm::FunctionCallee RegisterFunc = CGM.CreateRuntimeFunction(
363       llvm::FunctionType::get(IntTy, RegisterFuncParams, false),
364       addUnderscoredPrefixToName("RegisterFunction"));
365 
366   // Extract GpuBinaryHandle passed as the first argument passed to
367   // __cuda_register_globals() and generate __cudaRegisterFunction() call for
368   // each emitted kernel.
369   llvm::Argument &GpuBinaryHandlePtr = *RegisterKernelsFunc->arg_begin();
370   for (llvm::Function *Kernel : EmittedKernels) {
371     llvm::Constant *KernelName = makeConstantString(Kernel->getName());
372     llvm::Constant *NullPtr = llvm::ConstantPointerNull::get(VoidPtrTy);
373     llvm::Value *Args[] = {
374         &GpuBinaryHandlePtr, Builder.CreateBitCast(Kernel, VoidPtrTy),
375         KernelName, KernelName, llvm::ConstantInt::get(IntTy, -1), NullPtr,
376         NullPtr, NullPtr, NullPtr,
377         llvm::ConstantPointerNull::get(IntTy->getPointerTo())};
378     Builder.CreateCall(RegisterFunc, Args);
379   }
380 
381   // void __cudaRegisterVar(void **, char *, char *, const char *,
382   //                        int, int, int, int)
383   llvm::Type *RegisterVarParams[] = {VoidPtrPtrTy, CharPtrTy, CharPtrTy,
384                                      CharPtrTy,    IntTy,     IntTy,
385                                      IntTy,        IntTy};
386   llvm::FunctionCallee RegisterVar = CGM.CreateRuntimeFunction(
387       llvm::FunctionType::get(IntTy, RegisterVarParams, false),
388       addUnderscoredPrefixToName("RegisterVar"));
389   for (auto &Pair : DeviceVars) {
390     llvm::GlobalVariable *Var = Pair.first;
391     unsigned Flags = Pair.second;
392     llvm::Constant *VarName = makeConstantString(Var->getName());
393     uint64_t VarSize =
394         CGM.getDataLayout().getTypeAllocSize(Var->getValueType());
395     llvm::Value *Args[] = {
396         &GpuBinaryHandlePtr,
397         Builder.CreateBitCast(Var, VoidPtrTy),
398         VarName,
399         VarName,
400         llvm::ConstantInt::get(IntTy, (Flags & ExternDeviceVar) ? 1 : 0),
401         llvm::ConstantInt::get(IntTy, VarSize),
402         llvm::ConstantInt::get(IntTy, (Flags & ConstantDeviceVar) ? 1 : 0),
403         llvm::ConstantInt::get(IntTy, 0)};
404     Builder.CreateCall(RegisterVar, Args);
405   }
406 
407   Builder.CreateRetVoid();
408   return RegisterKernelsFunc;
409 }
410 
411 /// Creates a global constructor function for the module:
412 ///
413 /// For CUDA:
414 /// \code
415 /// void __cuda_module_ctor(void*) {
416 ///     Handle = __cudaRegisterFatBinary(GpuBinaryBlob);
417 ///     __cuda_register_globals(Handle);
418 /// }
419 /// \endcode
420 ///
421 /// For HIP:
422 /// \code
423 /// void __hip_module_ctor(void*) {
424 ///     if (__hip_gpubin_handle == 0) {
425 ///         __hip_gpubin_handle  = __hipRegisterFatBinary(GpuBinaryBlob);
426 ///         __hip_register_globals(__hip_gpubin_handle);
427 ///     }
428 /// }
429 /// \endcode
430 llvm::Function *CGNVCUDARuntime::makeModuleCtorFunction() {
431   bool IsHIP = CGM.getLangOpts().HIP;
432   // No need to generate ctors/dtors if there is no GPU binary.
433   StringRef CudaGpuBinaryFileName = CGM.getCodeGenOpts().CudaGpuBinaryFileName;
434   if (CudaGpuBinaryFileName.empty() && !IsHIP)
435     return nullptr;
436 
437   // void __{cuda|hip}_register_globals(void* handle);
438   llvm::Function *RegisterGlobalsFunc = makeRegisterGlobalsFn();
439   // We always need a function to pass in as callback. Create a dummy
440   // implementation if we don't need to register anything.
441   if (RelocatableDeviceCode && !RegisterGlobalsFunc)
442     RegisterGlobalsFunc = makeDummyFunction(getRegisterGlobalsFnTy());
443 
444   // void ** __{cuda|hip}RegisterFatBinary(void *);
445   llvm::FunctionCallee RegisterFatbinFunc = CGM.CreateRuntimeFunction(
446       llvm::FunctionType::get(VoidPtrPtrTy, VoidPtrTy, false),
447       addUnderscoredPrefixToName("RegisterFatBinary"));
448   // struct { int magic, int version, void * gpu_binary, void * dont_care };
449   llvm::StructType *FatbinWrapperTy =
450       llvm::StructType::get(IntTy, IntTy, VoidPtrTy, VoidPtrTy);
451 
452   // Register GPU binary with the CUDA runtime, store returned handle in a
453   // global variable and save a reference in GpuBinaryHandle to be cleaned up
454   // in destructor on exit. Then associate all known kernels with the GPU binary
455   // handle so CUDA runtime can figure out what to call on the GPU side.
456   std::unique_ptr<llvm::MemoryBuffer> CudaGpuBinary = nullptr;
457   if (!CudaGpuBinaryFileName.empty()) {
458     llvm::ErrorOr<std::unique_ptr<llvm::MemoryBuffer>> CudaGpuBinaryOrErr =
459         llvm::MemoryBuffer::getFileOrSTDIN(CudaGpuBinaryFileName);
460     if (std::error_code EC = CudaGpuBinaryOrErr.getError()) {
461       CGM.getDiags().Report(diag::err_cannot_open_file)
462           << CudaGpuBinaryFileName << EC.message();
463       return nullptr;
464     }
465     CudaGpuBinary = std::move(CudaGpuBinaryOrErr.get());
466   }
467 
468   llvm::Function *ModuleCtorFunc = llvm::Function::Create(
469       llvm::FunctionType::get(VoidTy, VoidPtrTy, false),
470       llvm::GlobalValue::InternalLinkage,
471       addUnderscoredPrefixToName("_module_ctor"), &TheModule);
472   llvm::BasicBlock *CtorEntryBB =
473       llvm::BasicBlock::Create(Context, "entry", ModuleCtorFunc);
474   CGBuilderTy CtorBuilder(CGM, Context);
475 
476   CtorBuilder.SetInsertPoint(CtorEntryBB);
477 
478   const char *FatbinConstantName;
479   const char *FatbinSectionName;
480   const char *ModuleIDSectionName;
481   StringRef ModuleIDPrefix;
482   llvm::Constant *FatBinStr;
483   unsigned FatMagic;
484   if (IsHIP) {
485     FatbinConstantName = ".hip_fatbin";
486     FatbinSectionName = ".hipFatBinSegment";
487 
488     ModuleIDSectionName = "__hip_module_id";
489     ModuleIDPrefix = "__hip_";
490 
491     if (CudaGpuBinary) {
492       // If fatbin is available from early finalization, create a string
493       // literal containing the fat binary loaded from the given file.
494       FatBinStr = makeConstantString(CudaGpuBinary->getBuffer(), "",
495                                      FatbinConstantName, 8);
496     } else {
497       // If fatbin is not available, create an external symbol
498       // __hip_fatbin in section .hip_fatbin. The external symbol is supposed
499       // to contain the fat binary but will be populated somewhere else,
500       // e.g. by lld through link script.
501       FatBinStr = new llvm::GlobalVariable(
502         CGM.getModule(), CGM.Int8Ty,
503         /*isConstant=*/true, llvm::GlobalValue::ExternalLinkage, nullptr,
504         "__hip_fatbin", nullptr,
505         llvm::GlobalVariable::NotThreadLocal);
506       cast<llvm::GlobalVariable>(FatBinStr)->setSection(FatbinConstantName);
507     }
508 
509     FatMagic = HIPFatMagic;
510   } else {
511     if (RelocatableDeviceCode)
512       FatbinConstantName = CGM.getTriple().isMacOSX()
513                                ? "__NV_CUDA,__nv_relfatbin"
514                                : "__nv_relfatbin";
515     else
516       FatbinConstantName =
517           CGM.getTriple().isMacOSX() ? "__NV_CUDA,__nv_fatbin" : ".nv_fatbin";
518     // NVIDIA's cuobjdump looks for fatbins in this section.
519     FatbinSectionName =
520         CGM.getTriple().isMacOSX() ? "__NV_CUDA,__fatbin" : ".nvFatBinSegment";
521 
522     ModuleIDSectionName = CGM.getTriple().isMacOSX()
523                               ? "__NV_CUDA,__nv_module_id"
524                               : "__nv_module_id";
525     ModuleIDPrefix = "__nv_";
526 
527     // For CUDA, create a string literal containing the fat binary loaded from
528     // the given file.
529     FatBinStr = makeConstantString(CudaGpuBinary->getBuffer(), "",
530                                    FatbinConstantName, 8);
531     FatMagic = CudaFatMagic;
532   }
533 
534   // Create initialized wrapper structure that points to the loaded GPU binary
535   ConstantInitBuilder Builder(CGM);
536   auto Values = Builder.beginStruct(FatbinWrapperTy);
537   // Fatbin wrapper magic.
538   Values.addInt(IntTy, FatMagic);
539   // Fatbin version.
540   Values.addInt(IntTy, 1);
541   // Data.
542   Values.add(FatBinStr);
543   // Unused in fatbin v1.
544   Values.add(llvm::ConstantPointerNull::get(VoidPtrTy));
545   llvm::GlobalVariable *FatbinWrapper = Values.finishAndCreateGlobal(
546       addUnderscoredPrefixToName("_fatbin_wrapper"), CGM.getPointerAlign(),
547       /*constant*/ true);
548   FatbinWrapper->setSection(FatbinSectionName);
549 
550   // There is only one HIP fat binary per linked module, however there are
551   // multiple constructor functions. Make sure the fat binary is registered
552   // only once. The constructor functions are executed by the dynamic loader
553   // before the program gains control. The dynamic loader cannot execute the
554   // constructor functions concurrently since doing that would not guarantee
555   // thread safety of the loaded program. Therefore we can assume sequential
556   // execution of constructor functions here.
557   if (IsHIP) {
558     auto Linkage = CudaGpuBinary ? llvm::GlobalValue::InternalLinkage :
559         llvm::GlobalValue::LinkOnceAnyLinkage;
560     llvm::BasicBlock *IfBlock =
561         llvm::BasicBlock::Create(Context, "if", ModuleCtorFunc);
562     llvm::BasicBlock *ExitBlock =
563         llvm::BasicBlock::Create(Context, "exit", ModuleCtorFunc);
564     // The name, size, and initialization pattern of this variable is part
565     // of HIP ABI.
566     GpuBinaryHandle = new llvm::GlobalVariable(
567         TheModule, VoidPtrPtrTy, /*isConstant=*/false,
568         Linkage,
569         /*Initializer=*/llvm::ConstantPointerNull::get(VoidPtrPtrTy),
570         "__hip_gpubin_handle");
571     GpuBinaryHandle->setAlignment(CGM.getPointerAlign().getQuantity());
572     // Prevent the weak symbol in different shared libraries being merged.
573     if (Linkage != llvm::GlobalValue::InternalLinkage)
574       GpuBinaryHandle->setVisibility(llvm::GlobalValue::HiddenVisibility);
575     Address GpuBinaryAddr(
576         GpuBinaryHandle,
577         CharUnits::fromQuantity(GpuBinaryHandle->getAlignment()));
578     {
579       auto HandleValue = CtorBuilder.CreateLoad(GpuBinaryAddr);
580       llvm::Constant *Zero =
581           llvm::Constant::getNullValue(HandleValue->getType());
582       llvm::Value *EQZero = CtorBuilder.CreateICmpEQ(HandleValue, Zero);
583       CtorBuilder.CreateCondBr(EQZero, IfBlock, ExitBlock);
584     }
585     {
586       CtorBuilder.SetInsertPoint(IfBlock);
587       // GpuBinaryHandle = __hipRegisterFatBinary(&FatbinWrapper);
588       llvm::CallInst *RegisterFatbinCall = CtorBuilder.CreateCall(
589           RegisterFatbinFunc,
590           CtorBuilder.CreateBitCast(FatbinWrapper, VoidPtrTy));
591       CtorBuilder.CreateStore(RegisterFatbinCall, GpuBinaryAddr);
592       CtorBuilder.CreateBr(ExitBlock);
593     }
594     {
595       CtorBuilder.SetInsertPoint(ExitBlock);
596       // Call __hip_register_globals(GpuBinaryHandle);
597       if (RegisterGlobalsFunc) {
598         auto HandleValue = CtorBuilder.CreateLoad(GpuBinaryAddr);
599         CtorBuilder.CreateCall(RegisterGlobalsFunc, HandleValue);
600       }
601     }
602   } else if (!RelocatableDeviceCode) {
603     // Register binary with CUDA runtime. This is substantially different in
604     // default mode vs. separate compilation!
605     // GpuBinaryHandle = __cudaRegisterFatBinary(&FatbinWrapper);
606     llvm::CallInst *RegisterFatbinCall = CtorBuilder.CreateCall(
607         RegisterFatbinFunc,
608         CtorBuilder.CreateBitCast(FatbinWrapper, VoidPtrTy));
609     GpuBinaryHandle = new llvm::GlobalVariable(
610         TheModule, VoidPtrPtrTy, false, llvm::GlobalValue::InternalLinkage,
611         llvm::ConstantPointerNull::get(VoidPtrPtrTy), "__cuda_gpubin_handle");
612     GpuBinaryHandle->setAlignment(CGM.getPointerAlign().getQuantity());
613     CtorBuilder.CreateAlignedStore(RegisterFatbinCall, GpuBinaryHandle,
614                                    CGM.getPointerAlign());
615 
616     // Call __cuda_register_globals(GpuBinaryHandle);
617     if (RegisterGlobalsFunc)
618       CtorBuilder.CreateCall(RegisterGlobalsFunc, RegisterFatbinCall);
619 
620     // Call __cudaRegisterFatBinaryEnd(Handle) if this CUDA version needs it.
621     if (CudaFeatureEnabled(CGM.getTarget().getSDKVersion(),
622                            CudaFeature::CUDA_USES_FATBIN_REGISTER_END)) {
623       // void __cudaRegisterFatBinaryEnd(void **);
624       llvm::FunctionCallee RegisterFatbinEndFunc = CGM.CreateRuntimeFunction(
625           llvm::FunctionType::get(VoidTy, VoidPtrPtrTy, false),
626           "__cudaRegisterFatBinaryEnd");
627       CtorBuilder.CreateCall(RegisterFatbinEndFunc, RegisterFatbinCall);
628     }
629   } else {
630     // Generate a unique module ID.
631     SmallString<64> ModuleID;
632     llvm::raw_svector_ostream OS(ModuleID);
633     OS << ModuleIDPrefix << llvm::format("%" PRIx64, FatbinWrapper->getGUID());
634     llvm::Constant *ModuleIDConstant =
635         makeConstantString(ModuleID.str(), "", ModuleIDSectionName, 32);
636 
637     // Create an alias for the FatbinWrapper that nvcc will look for.
638     llvm::GlobalAlias::create(llvm::GlobalValue::ExternalLinkage,
639                               Twine("__fatbinwrap") + ModuleID, FatbinWrapper);
640 
641     // void __cudaRegisterLinkedBinary%ModuleID%(void (*)(void *), void *,
642     // void *, void (*)(void **))
643     SmallString<128> RegisterLinkedBinaryName("__cudaRegisterLinkedBinary");
644     RegisterLinkedBinaryName += ModuleID;
645     llvm::FunctionCallee RegisterLinkedBinaryFunc = CGM.CreateRuntimeFunction(
646         getRegisterLinkedBinaryFnTy(), RegisterLinkedBinaryName);
647 
648     assert(RegisterGlobalsFunc && "Expecting at least dummy function!");
649     llvm::Value *Args[] = {RegisterGlobalsFunc,
650                            CtorBuilder.CreateBitCast(FatbinWrapper, VoidPtrTy),
651                            ModuleIDConstant,
652                            makeDummyFunction(getCallbackFnTy())};
653     CtorBuilder.CreateCall(RegisterLinkedBinaryFunc, Args);
654   }
655 
656   // Create destructor and register it with atexit() the way NVCC does it. Doing
657   // it during regular destructor phase worked in CUDA before 9.2 but results in
658   // double-free in 9.2.
659   if (llvm::Function *CleanupFn = makeModuleDtorFunction()) {
660     // extern "C" int atexit(void (*f)(void));
661     llvm::FunctionType *AtExitTy =
662         llvm::FunctionType::get(IntTy, CleanupFn->getType(), false);
663     llvm::FunctionCallee AtExitFunc =
664         CGM.CreateRuntimeFunction(AtExitTy, "atexit", llvm::AttributeList(),
665                                   /*Local=*/true);
666     CtorBuilder.CreateCall(AtExitFunc, CleanupFn);
667   }
668 
669   CtorBuilder.CreateRetVoid();
670   return ModuleCtorFunc;
671 }
672 
673 /// Creates a global destructor function that unregisters the GPU code blob
674 /// registered by constructor.
675 ///
676 /// For CUDA:
677 /// \code
678 /// void __cuda_module_dtor(void*) {
679 ///     __cudaUnregisterFatBinary(Handle);
680 /// }
681 /// \endcode
682 ///
683 /// For HIP:
684 /// \code
685 /// void __hip_module_dtor(void*) {
686 ///     if (__hip_gpubin_handle) {
687 ///         __hipUnregisterFatBinary(__hip_gpubin_handle);
688 ///         __hip_gpubin_handle = 0;
689 ///     }
690 /// }
691 /// \endcode
692 llvm::Function *CGNVCUDARuntime::makeModuleDtorFunction() {
693   // No need for destructor if we don't have a handle to unregister.
694   if (!GpuBinaryHandle)
695     return nullptr;
696 
697   // void __cudaUnregisterFatBinary(void ** handle);
698   llvm::FunctionCallee UnregisterFatbinFunc = CGM.CreateRuntimeFunction(
699       llvm::FunctionType::get(VoidTy, VoidPtrPtrTy, false),
700       addUnderscoredPrefixToName("UnregisterFatBinary"));
701 
702   llvm::Function *ModuleDtorFunc = llvm::Function::Create(
703       llvm::FunctionType::get(VoidTy, VoidPtrTy, false),
704       llvm::GlobalValue::InternalLinkage,
705       addUnderscoredPrefixToName("_module_dtor"), &TheModule);
706 
707   llvm::BasicBlock *DtorEntryBB =
708       llvm::BasicBlock::Create(Context, "entry", ModuleDtorFunc);
709   CGBuilderTy DtorBuilder(CGM, Context);
710   DtorBuilder.SetInsertPoint(DtorEntryBB);
711 
712   Address GpuBinaryAddr(GpuBinaryHandle, CharUnits::fromQuantity(
713                                              GpuBinaryHandle->getAlignment()));
714   auto HandleValue = DtorBuilder.CreateLoad(GpuBinaryAddr);
715   // There is only one HIP fat binary per linked module, however there are
716   // multiple destructor functions. Make sure the fat binary is unregistered
717   // only once.
718   if (CGM.getLangOpts().HIP) {
719     llvm::BasicBlock *IfBlock =
720         llvm::BasicBlock::Create(Context, "if", ModuleDtorFunc);
721     llvm::BasicBlock *ExitBlock =
722         llvm::BasicBlock::Create(Context, "exit", ModuleDtorFunc);
723     llvm::Constant *Zero = llvm::Constant::getNullValue(HandleValue->getType());
724     llvm::Value *NEZero = DtorBuilder.CreateICmpNE(HandleValue, Zero);
725     DtorBuilder.CreateCondBr(NEZero, IfBlock, ExitBlock);
726 
727     DtorBuilder.SetInsertPoint(IfBlock);
728     DtorBuilder.CreateCall(UnregisterFatbinFunc, HandleValue);
729     DtorBuilder.CreateStore(Zero, GpuBinaryAddr);
730     DtorBuilder.CreateBr(ExitBlock);
731 
732     DtorBuilder.SetInsertPoint(ExitBlock);
733   } else {
734     DtorBuilder.CreateCall(UnregisterFatbinFunc, HandleValue);
735   }
736   DtorBuilder.CreateRetVoid();
737   return ModuleDtorFunc;
738 }
739 
740 CGCUDARuntime *CodeGen::CreateNVCUDARuntime(CodeGenModule &CGM) {
741   return new CGNVCUDARuntime(CGM);
742 }
743