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