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