clang  15.0.0git
SemaCUDA.cpp
Go to the documentation of this file.
1 //===--- SemaCUDA.cpp - Semantic Analysis for CUDA constructs -------------===//
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 /// \file
9 /// This file implements semantic analysis for CUDA constructs.
10 ///
11 //===----------------------------------------------------------------------===//
12 
13 #include "clang/AST/ASTContext.h"
14 #include "clang/AST/Decl.h"
15 #include "clang/AST/ExprCXX.h"
16 #include "clang/Basic/Cuda.h"
17 #include "clang/Basic/TargetInfo.h"
18 #include "clang/Lex/Preprocessor.h"
19 #include "clang/Sema/Lookup.h"
20 #include "clang/Sema/ScopeInfo.h"
21 #include "clang/Sema/Sema.h"
24 #include "clang/Sema/Template.h"
25 #include "llvm/ADT/Optional.h"
26 #include "llvm/ADT/SmallVector.h"
27 using namespace clang;
28 
29 template <typename AttrT> static bool hasExplicitAttr(const VarDecl *D) {
30  if (!D)
31  return false;
32  if (auto *A = D->getAttr<AttrT>())
33  return !A->isImplicit();
34  return false;
35 }
36 
38  assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
39  ForceCUDAHostDeviceDepth++;
40 }
41 
43  assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
44  if (ForceCUDAHostDeviceDepth == 0)
45  return false;
46  ForceCUDAHostDeviceDepth--;
47  return true;
48 }
49 
51  MultiExprArg ExecConfig,
52  SourceLocation GGGLoc) {
54  if (!ConfigDecl)
55  return ExprError(Diag(LLLLoc, diag::err_undeclared_var_use)
57  QualType ConfigQTy = ConfigDecl->getType();
58 
59  DeclRefExpr *ConfigDR = new (Context)
60  DeclRefExpr(Context, ConfigDecl, false, ConfigQTy, VK_LValue, LLLLoc);
61  MarkFunctionReferenced(LLLLoc, ConfigDecl);
62 
63  return BuildCallExpr(S, ConfigDR, LLLLoc, ExecConfig, GGGLoc, nullptr,
64  /*IsExecConfig=*/true);
65 }
66 
69  bool HasHostAttr = false;
70  bool HasDeviceAttr = false;
71  bool HasGlobalAttr = false;
72  bool HasInvalidTargetAttr = false;
73  for (const ParsedAttr &AL : Attrs) {
74  switch (AL.getKind()) {
75  case ParsedAttr::AT_CUDAGlobal:
76  HasGlobalAttr = true;
77  break;
78  case ParsedAttr::AT_CUDAHost:
79  HasHostAttr = true;
80  break;
81  case ParsedAttr::AT_CUDADevice:
82  HasDeviceAttr = true;
83  break;
84  case ParsedAttr::AT_CUDAInvalidTarget:
85  HasInvalidTargetAttr = true;
86  break;
87  default:
88  break;
89  }
90  }
91 
92  if (HasInvalidTargetAttr)
93  return CFT_InvalidTarget;
94 
95  if (HasGlobalAttr)
96  return CFT_Global;
97 
98  if (HasHostAttr && HasDeviceAttr)
99  return CFT_HostDevice;
100 
101  if (HasDeviceAttr)
102  return CFT_Device;
103 
104  return CFT_Host;
105 }
106 
107 template <typename A>
108 static bool hasAttr(const FunctionDecl *D, bool IgnoreImplicitAttr) {
109  return D->hasAttrs() && llvm::any_of(D->getAttrs(), [&](Attr *Attribute) {
110  return isa<A>(Attribute) &&
111  !(IgnoreImplicitAttr && Attribute->isImplicit());
112  });
113 }
114 
115 /// IdentifyCUDATarget - Determine the CUDA compilation target for this function
117  bool IgnoreImplicitHDAttr) {
118  // Code that lives outside a function is run on the host.
119  if (D == nullptr)
120  return CFT_Host;
121 
122  if (D->hasAttr<CUDAInvalidTargetAttr>())
123  return CFT_InvalidTarget;
124 
125  if (D->hasAttr<CUDAGlobalAttr>())
126  return CFT_Global;
127 
128  if (hasAttr<CUDADeviceAttr>(D, IgnoreImplicitHDAttr)) {
129  if (hasAttr<CUDAHostAttr>(D, IgnoreImplicitHDAttr))
130  return CFT_HostDevice;
131  return CFT_Device;
132  } else if (hasAttr<CUDAHostAttr>(D, IgnoreImplicitHDAttr)) {
133  return CFT_Host;
134  } else if ((D->isImplicit() || !D->isUserProvided()) &&
135  !IgnoreImplicitHDAttr) {
136  // Some implicit declarations (like intrinsic functions) are not marked.
137  // Set the most lenient target on them for maximal flexibility.
138  return CFT_HostDevice;
139  }
140 
141  return CFT_Host;
142 }
143 
144 /// IdentifyTarget - Determine the CUDA compilation target for this variable.
146  if (Var->hasAttr<HIPManagedAttr>())
147  return CVT_Unified;
148  // Only constexpr and const variabless with implicit constant attribute
149  // are emitted on both sides. Such variables are promoted to device side
150  // only if they have static constant intializers on device side.
151  if ((Var->isConstexpr() || Var->getType().isConstQualified()) &&
152  Var->hasAttr<CUDAConstantAttr>() &&
153  !hasExplicitAttr<CUDAConstantAttr>(Var))
154  return CVT_Both;
155  if (Var->hasAttr<CUDADeviceAttr>() || Var->hasAttr<CUDAConstantAttr>() ||
156  Var->hasAttr<CUDASharedAttr>() ||
159  return CVT_Device;
160  // Function-scope static variable without explicit device or constant
161  // attribute are emitted
162  // - on both sides in host device functions
163  // - on device side in device or global functions
164  if (auto *FD = dyn_cast<FunctionDecl>(Var->getDeclContext())) {
165  switch (IdentifyCUDATarget(FD)) {
166  case CFT_HostDevice:
167  return CVT_Both;
168  case CFT_Device:
169  case CFT_Global:
170  return CVT_Device;
171  default:
172  return CVT_Host;
173  }
174  }
175  return CVT_Host;
176 }
177 
178 // * CUDA Call preference table
179 //
180 // F - from,
181 // T - to
182 // Ph - preference in host mode
183 // Pd - preference in device mode
184 // H - handled in (x)
185 // Preferences: N:native, SS:same side, HD:host-device, WS:wrong side, --:never.
186 //
187 // | F | T | Ph | Pd | H |
188 // |----+----+-----+-----+-----+
189 // | d | d | N | N | (c) |
190 // | d | g | -- | -- | (a) |
191 // | d | h | -- | -- | (e) |
192 // | d | hd | HD | HD | (b) |
193 // | g | d | N | N | (c) |
194 // | g | g | -- | -- | (a) |
195 // | g | h | -- | -- | (e) |
196 // | g | hd | HD | HD | (b) |
197 // | h | d | -- | -- | (e) |
198 // | h | g | N | N | (c) |
199 // | h | h | N | N | (c) |
200 // | h | hd | HD | HD | (b) |
201 // | hd | d | WS | SS | (d) |
202 // | hd | g | SS | -- |(d/a)|
203 // | hd | h | SS | WS | (d) |
204 // | hd | hd | HD | HD | (b) |
205 
208  const FunctionDecl *Callee) {
209  assert(Callee && "Callee must be valid.");
210  CUDAFunctionTarget CallerTarget = IdentifyCUDATarget(Caller);
211  CUDAFunctionTarget CalleeTarget = IdentifyCUDATarget(Callee);
212 
213  // If one of the targets is invalid, the check always fails, no matter what
214  // the other target is.
215  if (CallerTarget == CFT_InvalidTarget || CalleeTarget == CFT_InvalidTarget)
216  return CFP_Never;
217 
218  // (a) Can't call global from some contexts until we support CUDA's
219  // dynamic parallelism.
220  if (CalleeTarget == CFT_Global &&
221  (CallerTarget == CFT_Global || CallerTarget == CFT_Device))
222  return CFP_Never;
223 
224  // (b) Calling HostDevice is OK for everyone.
225  if (CalleeTarget == CFT_HostDevice)
226  return CFP_HostDevice;
227 
228  // (c) Best case scenarios
229  if (CalleeTarget == CallerTarget ||
230  (CallerTarget == CFT_Host && CalleeTarget == CFT_Global) ||
231  (CallerTarget == CFT_Global && CalleeTarget == CFT_Device))
232  return CFP_Native;
233 
234  // (d) HostDevice behavior depends on compilation mode.
235  if (CallerTarget == CFT_HostDevice) {
236  // It's OK to call a compilation-mode matching function from an HD one.
237  if ((getLangOpts().CUDAIsDevice && CalleeTarget == CFT_Device) ||
238  (!getLangOpts().CUDAIsDevice &&
239  (CalleeTarget == CFT_Host || CalleeTarget == CFT_Global)))
240  return CFP_SameSide;
241 
242  // Calls from HD to non-mode-matching functions (i.e., to host functions
243  // when compiling in device mode or to device functions when compiling in
244  // host mode) are allowed at the sema level, but eventually rejected if
245  // they're ever codegened. TODO: Reject said calls earlier.
246  return CFP_WrongSide;
247  }
248 
249  // (e) Calling across device/host boundary is not something you should do.
250  if ((CallerTarget == CFT_Host && CalleeTarget == CFT_Device) ||
251  (CallerTarget == CFT_Device && CalleeTarget == CFT_Host) ||
252  (CallerTarget == CFT_Global && CalleeTarget == CFT_Host))
253  return CFP_Never;
254 
255  llvm_unreachable("All cases should've been handled by now.");
256 }
257 
258 template <typename AttrT> static bool hasImplicitAttr(const FunctionDecl *D) {
259  if (!D)
260  return false;
261  if (auto *A = D->getAttr<AttrT>())
262  return A->isImplicit();
263  return D->isImplicit();
264 }
265 
267  bool IsImplicitDevAttr = hasImplicitAttr<CUDADeviceAttr>(D);
268  bool IsImplicitHostAttr = hasImplicitAttr<CUDAHostAttr>(D);
269  return IsImplicitDevAttr && IsImplicitHostAttr;
270 }
271 
273  const FunctionDecl *Caller,
274  SmallVectorImpl<std::pair<DeclAccessPair, FunctionDecl *>> &Matches) {
275  if (Matches.size() <= 1)
276  return;
277 
278  using Pair = std::pair<DeclAccessPair, FunctionDecl*>;
279 
280  // Gets the CUDA function preference for a call from Caller to Match.
281  auto GetCFP = [&](const Pair &Match) {
282  return IdentifyCUDAPreference(Caller, Match.second);
283  };
284 
285  // Find the best call preference among the functions in Matches.
286  CUDAFunctionPreference BestCFP = GetCFP(*std::max_element(
287  Matches.begin(), Matches.end(),
288  [&](const Pair &M1, const Pair &M2) { return GetCFP(M1) < GetCFP(M2); }));
289 
290  // Erase all functions with lower priority.
291  llvm::erase_if(Matches,
292  [&](const Pair &Match) { return GetCFP(Match) < BestCFP; });
293 }
294 
295 /// When an implicitly-declared special member has to invoke more than one
296 /// base/field special member, conflicts may occur in the targets of these
297 /// members. For example, if one base's member __host__ and another's is
298 /// __device__, it's a conflict.
299 /// This function figures out if the given targets \param Target1 and
300 /// \param Target2 conflict, and if they do not it fills in
301 /// \param ResolvedTarget with a target that resolves for both calls.
302 /// \return true if there's a conflict, false otherwise.
303 static bool
305  Sema::CUDAFunctionTarget Target2,
306  Sema::CUDAFunctionTarget *ResolvedTarget) {
307  // Only free functions and static member functions may be global.
308  assert(Target1 != Sema::CFT_Global);
309  assert(Target2 != Sema::CFT_Global);
310 
311  if (Target1 == Sema::CFT_HostDevice) {
312  *ResolvedTarget = Target2;
313  } else if (Target2 == Sema::CFT_HostDevice) {
314  *ResolvedTarget = Target1;
315  } else if (Target1 != Target2) {
316  return true;
317  } else {
318  *ResolvedTarget = Target1;
319  }
320 
321  return false;
322 }
323 
325  CXXSpecialMember CSM,
326  CXXMethodDecl *MemberDecl,
327  bool ConstRHS,
328  bool Diagnose) {
329  // If the defaulted special member is defined lexically outside of its
330  // owning class, or the special member already has explicit device or host
331  // attributes, do not infer.
332  bool InClass = MemberDecl->getLexicalParent() == MemberDecl->getParent();
333  bool HasH = MemberDecl->hasAttr<CUDAHostAttr>();
334  bool HasD = MemberDecl->hasAttr<CUDADeviceAttr>();
335  bool HasExplicitAttr =
336  (HasD && !MemberDecl->getAttr<CUDADeviceAttr>()->isImplicit()) ||
337  (HasH && !MemberDecl->getAttr<CUDAHostAttr>()->isImplicit());
338  if (!InClass || HasExplicitAttr)
339  return false;
340 
341  llvm::Optional<CUDAFunctionTarget> InferredTarget;
342 
343  // We're going to invoke special member lookup; mark that these special
344  // members are called from this one, and not from its caller.
345  ContextRAII MethodContext(*this, MemberDecl);
346 
347  // Look for special members in base classes that should be invoked from here.
348  // Infer the target of this member base on the ones it should call.
349  // Skip direct and indirect virtual bases for abstract classes.
351  for (const auto &B : ClassDecl->bases()) {
352  if (!B.isVirtual()) {
353  Bases.push_back(&B);
354  }
355  }
356 
357  if (!ClassDecl->isAbstract()) {
358  llvm::append_range(Bases, llvm::make_pointer_range(ClassDecl->vbases()));
359  }
360 
361  for (const auto *B : Bases) {
362  const RecordType *BaseType = B->getType()->getAs<RecordType>();
363  if (!BaseType) {
364  continue;
365  }
366 
367  CXXRecordDecl *BaseClassDecl = cast<CXXRecordDecl>(BaseType->getDecl());
369  LookupSpecialMember(BaseClassDecl, CSM,
370  /* ConstArg */ ConstRHS,
371  /* VolatileArg */ false,
372  /* RValueThis */ false,
373  /* ConstThis */ false,
374  /* VolatileThis */ false);
375 
376  if (!SMOR.getMethod())
377  continue;
378 
379  CUDAFunctionTarget BaseMethodTarget = IdentifyCUDATarget(SMOR.getMethod());
380  if (!InferredTarget.hasValue()) {
381  InferredTarget = BaseMethodTarget;
382  } else {
383  bool ResolutionError = resolveCalleeCUDATargetConflict(
384  InferredTarget.getValue(), BaseMethodTarget,
385  InferredTarget.getPointer());
386  if (ResolutionError) {
387  if (Diagnose) {
388  Diag(ClassDecl->getLocation(),
389  diag::note_implicit_member_target_infer_collision)
390  << (unsigned)CSM << InferredTarget.getValue() << BaseMethodTarget;
391  }
392  MemberDecl->addAttr(CUDAInvalidTargetAttr::CreateImplicit(Context));
393  return true;
394  }
395  }
396  }
397 
398  // Same as for bases, but now for special members of fields.
399  for (const auto *F : ClassDecl->fields()) {
400  if (F->isInvalidDecl()) {
401  continue;
402  }
403 
404  const RecordType *FieldType =
405  Context.getBaseElementType(F->getType())->getAs<RecordType>();
406  if (!FieldType) {
407  continue;
408  }
409 
410  CXXRecordDecl *FieldRecDecl = cast<CXXRecordDecl>(FieldType->getDecl());
412  LookupSpecialMember(FieldRecDecl, CSM,
413  /* ConstArg */ ConstRHS && !F->isMutable(),
414  /* VolatileArg */ false,
415  /* RValueThis */ false,
416  /* ConstThis */ false,
417  /* VolatileThis */ false);
418 
419  if (!SMOR.getMethod())
420  continue;
421 
422  CUDAFunctionTarget FieldMethodTarget =
424  if (!InferredTarget.hasValue()) {
425  InferredTarget = FieldMethodTarget;
426  } else {
427  bool ResolutionError = resolveCalleeCUDATargetConflict(
428  InferredTarget.getValue(), FieldMethodTarget,
429  InferredTarget.getPointer());
430  if (ResolutionError) {
431  if (Diagnose) {
432  Diag(ClassDecl->getLocation(),
433  diag::note_implicit_member_target_infer_collision)
434  << (unsigned)CSM << InferredTarget.getValue()
435  << FieldMethodTarget;
436  }
437  MemberDecl->addAttr(CUDAInvalidTargetAttr::CreateImplicit(Context));
438  return true;
439  }
440  }
441  }
442 
443 
444  // If no target was inferred, mark this member as __host__ __device__;
445  // it's the least restrictive option that can be invoked from any target.
446  bool NeedsH = true, NeedsD = true;
447  if (InferredTarget.hasValue()) {
448  if (InferredTarget.getValue() == CFT_Device)
449  NeedsH = false;
450  else if (InferredTarget.getValue() == CFT_Host)
451  NeedsD = false;
452  }
453 
454  // We either setting attributes first time, or the inferred ones must match
455  // previously set ones.
456  if (NeedsD && !HasD)
457  MemberDecl->addAttr(CUDADeviceAttr::CreateImplicit(Context));
458  if (NeedsH && !HasH)
459  MemberDecl->addAttr(CUDAHostAttr::CreateImplicit(Context));
460 
461  return false;
462 }
463 
465  if (!CD->isDefined() && CD->isTemplateInstantiation())
467 
468  // (E.2.3.1, CUDA 7.5) A constructor for a class type is considered
469  // empty at a point in the translation unit, if it is either a
470  // trivial constructor
471  if (CD->isTrivial())
472  return true;
473 
474  // ... or it satisfies all of the following conditions:
475  // The constructor function has been defined.
476  // The constructor function has no parameters,
477  // and the function body is an empty compound statement.
478  if (!(CD->hasTrivialBody() && CD->getNumParams() == 0))
479  return false;
480 
481  // Its class has no virtual functions and no virtual base classes.
482  if (CD->getParent()->isDynamicClass())
483  return false;
484 
485  // Union ctor does not call ctors of its data members.
486  if (CD->getParent()->isUnion())
487  return true;
488 
489  // The only form of initializer allowed is an empty constructor.
490  // This will recursively check all base classes and member initializers
491  if (!llvm::all_of(CD->inits(), [&](const CXXCtorInitializer *CI) {
492  if (const CXXConstructExpr *CE =
493  dyn_cast<CXXConstructExpr>(CI->getInit()))
494  return isEmptyCudaConstructor(Loc, CE->getConstructor());
495  return false;
496  }))
497  return false;
498 
499  return true;
500 }
501 
503  // No destructor -> no problem.
504  if (!DD)
505  return true;
506 
507  if (!DD->isDefined() && DD->isTemplateInstantiation())
509 
510  // (E.2.3.1, CUDA 7.5) A destructor for a class type is considered
511  // empty at a point in the translation unit, if it is either a
512  // trivial constructor
513  if (DD->isTrivial())
514  return true;
515 
516  // ... or it satisfies all of the following conditions:
517  // The destructor function has been defined.
518  // and the function body is an empty compound statement.
519  if (!DD->hasTrivialBody())
520  return false;
521 
522  const CXXRecordDecl *ClassDecl = DD->getParent();
523 
524  // Its class has no virtual functions and no virtual base classes.
525  if (ClassDecl->isDynamicClass())
526  return false;
527 
528  // Union does not have base class and union dtor does not call dtors of its
529  // data members.
530  if (DD->getParent()->isUnion())
531  return true;
532 
533  // Only empty destructors are allowed. This will recursively check
534  // destructors for all base classes...
535  if (!llvm::all_of(ClassDecl->bases(), [&](const CXXBaseSpecifier &BS) {
536  if (CXXRecordDecl *RD = BS.getType()->getAsCXXRecordDecl())
537  return isEmptyCudaDestructor(Loc, RD->getDestructor());
538  return true;
539  }))
540  return false;
541 
542  // ... and member fields.
543  if (!llvm::all_of(ClassDecl->fields(), [&](const FieldDecl *Field) {
544  if (CXXRecordDecl *RD = Field->getType()
545  ->getBaseElementTypeUnsafe()
546  ->getAsCXXRecordDecl())
547  return isEmptyCudaDestructor(Loc, RD->getDestructor());
548  return true;
549  }))
550  return false;
551 
552  return true;
553 }
554 
555 namespace {
556 enum CUDAInitializerCheckKind {
557  CICK_DeviceOrConstant, // Check initializer for device/constant variable
558  CICK_Shared, // Check initializer for shared variable
559 };
560 
561 bool IsDependentVar(VarDecl *VD) {
562  if (VD->getType()->isDependentType())
563  return true;
564  if (const auto *Init = VD->getInit())
565  return Init->isValueDependent();
566  return false;
567 }
568 
569 // Check whether a variable has an allowed initializer for a CUDA device side
570 // variable with global storage. \p VD may be a host variable to be checked for
571 // potential promotion to device side variable.
572 //
573 // CUDA/HIP allows only empty constructors as initializers for global
574 // variables (see E.2.3.1, CUDA 7.5). The same restriction also applies to all
575 // __shared__ variables whether they are local or not (they all are implicitly
576 // static in CUDA). One exception is that CUDA allows constant initializers
577 // for __constant__ and __device__ variables.
578 bool HasAllowedCUDADeviceStaticInitializer(Sema &S, VarDecl *VD,
579  CUDAInitializerCheckKind CheckKind) {
580  assert(!VD->isInvalidDecl() && VD->hasGlobalStorage());
581  assert(!IsDependentVar(VD) && "do not check dependent var");
582  const Expr *Init = VD->getInit();
583  auto IsEmptyInit = [&](const Expr *Init) {
584  if (!Init)
585  return true;
586  if (const auto *CE = dyn_cast<CXXConstructExpr>(Init)) {
587  return S.isEmptyCudaConstructor(VD->getLocation(), CE->getConstructor());
588  }
589  return false;
590  };
591  auto IsConstantInit = [&](const Expr *Init) {
592  assert(Init);
594  /*NoWronSidedVars=*/true);
595  return Init->isConstantInitializer(S.Context,
596  VD->getType()->isReferenceType());
597  };
598  auto HasEmptyDtor = [&](VarDecl *VD) {
599  if (const auto *RD = VD->getType()->getAsCXXRecordDecl())
600  return S.isEmptyCudaDestructor(VD->getLocation(), RD->getDestructor());
601  return true;
602  };
603  if (CheckKind == CICK_Shared)
604  return IsEmptyInit(Init) && HasEmptyDtor(VD);
605  return S.LangOpts.GPUAllowDeviceInit ||
606  ((IsEmptyInit(Init) || IsConstantInit(Init)) && HasEmptyDtor(VD));
607 }
608 } // namespace
609 
611  // Do not check dependent variables since the ctor/dtor/initializer are not
612  // determined. Do it after instantiation.
613  if (VD->isInvalidDecl() || !VD->hasInit() || !VD->hasGlobalStorage() ||
614  IsDependentVar(VD))
615  return;
616  const Expr *Init = VD->getInit();
617  bool IsSharedVar = VD->hasAttr<CUDASharedAttr>();
618  bool IsDeviceOrConstantVar =
619  !IsSharedVar &&
620  (VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>());
621  if (IsDeviceOrConstantVar || IsSharedVar) {
622  if (HasAllowedCUDADeviceStaticInitializer(
623  *this, VD, IsSharedVar ? CICK_Shared : CICK_DeviceOrConstant))
624  return;
625  Diag(VD->getLocation(),
626  IsSharedVar ? diag::err_shared_var_init : diag::err_dynamic_var_init)
627  << Init->getSourceRange();
628  VD->setInvalidDecl();
629  } else {
630  // This is a host-side global variable. Check that the initializer is
631  // callable from the host side.
632  const FunctionDecl *InitFn = nullptr;
633  if (const CXXConstructExpr *CE = dyn_cast<CXXConstructExpr>(Init)) {
634  InitFn = CE->getConstructor();
635  } else if (const CallExpr *CE = dyn_cast<CallExpr>(Init)) {
636  InitFn = CE->getDirectCallee();
637  }
638  if (InitFn) {
639  CUDAFunctionTarget InitFnTarget = IdentifyCUDATarget(InitFn);
640  if (InitFnTarget != CFT_Host && InitFnTarget != CFT_HostDevice) {
641  Diag(VD->getLocation(), diag::err_ref_bad_target_global_initializer)
642  << InitFnTarget << InitFn;
643  Diag(InitFn->getLocation(), diag::note_previous_decl) << InitFn;
644  VD->setInvalidDecl();
645  }
646  }
647  }
648 }
649 
650 // With -fcuda-host-device-constexpr, an unattributed constexpr function is
651 // treated as implicitly __host__ __device__, unless:
652 // * it is a variadic function (device-side variadic functions are not
653 // allowed), or
654 // * a __device__ function with this signature was already declared, in which
655 // case in which case we output an error, unless the __device__ decl is in a
656 // system header, in which case we leave the constexpr function unattributed.
657 //
658 // In addition, all function decls are treated as __host__ __device__ when
659 // ForceCUDAHostDeviceDepth > 0 (corresponding to code within a
660 // #pragma clang force_cuda_host_device_begin/end
661 // pair).
663  const LookupResult &Previous) {
664  assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
665 
666  if (ForceCUDAHostDeviceDepth > 0) {
667  if (!NewD->hasAttr<CUDAHostAttr>())
668  NewD->addAttr(CUDAHostAttr::CreateImplicit(Context));
669  if (!NewD->hasAttr<CUDADeviceAttr>())
670  NewD->addAttr(CUDADeviceAttr::CreateImplicit(Context));
671  return;
672  }
673 
674  if (!getLangOpts().CUDAHostDeviceConstexpr || !NewD->isConstexpr() ||
675  NewD->isVariadic() || NewD->hasAttr<CUDAHostAttr>() ||
676  NewD->hasAttr<CUDADeviceAttr>() || NewD->hasAttr<CUDAGlobalAttr>())
677  return;
678 
679  // Is D a __device__ function with the same signature as NewD, ignoring CUDA
680  // attributes?
681  auto IsMatchingDeviceFn = [&](NamedDecl *D) {
682  if (UsingShadowDecl *Using = dyn_cast<UsingShadowDecl>(D))
683  D = Using->getTargetDecl();
684  FunctionDecl *OldD = D->getAsFunction();
685  return OldD && OldD->hasAttr<CUDADeviceAttr>() &&
686  !OldD->hasAttr<CUDAHostAttr>() &&
687  !IsOverload(NewD, OldD, /* UseMemberUsingDeclRules = */ false,
688  /* ConsiderCudaAttrs = */ false);
689  };
690  auto It = llvm::find_if(Previous, IsMatchingDeviceFn);
691  if (It != Previous.end()) {
692  // We found a __device__ function with the same name and signature as NewD
693  // (ignoring CUDA attrs). This is an error unless that function is defined
694  // in a system header, in which case we simply return without making NewD
695  // host+device.
696  NamedDecl *Match = *It;
697  if (!getSourceManager().isInSystemHeader(Match->getLocation())) {
698  Diag(NewD->getLocation(),
699  diag::err_cuda_unattributed_constexpr_cannot_overload_device)
700  << NewD;
701  Diag(Match->getLocation(),
702  diag::note_cuda_conflicting_device_function_declared_here);
703  }
704  return;
705  }
706 
707  NewD->addAttr(CUDAHostAttr::CreateImplicit(Context));
708  NewD->addAttr(CUDADeviceAttr::CreateImplicit(Context));
709 }
710 
711 // TODO: `__constant__` memory may be a limited resource for certain targets.
712 // A safeguard may be needed at the end of compilation pipeline if
713 // `__constant__` memory usage goes beyond limit.
715  // Do not promote dependent variables since the cotr/dtor/initializer are
716  // not determined. Do it after instantiation.
717  if (getLangOpts().CUDAIsDevice && !VD->hasAttr<CUDAConstantAttr>() &&
718  !VD->hasAttr<CUDAConstantAttr>() && !VD->hasAttr<CUDASharedAttr>() &&
719  (VD->isFileVarDecl() || VD->isStaticDataMember()) &&
720  !IsDependentVar(VD) &&
721  ((VD->isConstexpr() || VD->getType().isConstQualified()) &&
722  HasAllowedCUDADeviceStaticInitializer(*this, VD,
723  CICK_DeviceOrConstant))) {
724  VD->addAttr(CUDAConstantAttr::CreateImplicit(getASTContext()));
725  }
726 }
727 
729  unsigned DiagID) {
730  assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
731  FunctionDecl *CurFunContext = getCurFunctionDecl(/*AllowLambda=*/true);
732  SemaDiagnosticBuilder::Kind DiagKind = [&] {
733  if (!CurFunContext)
735  switch (CurrentCUDATarget()) {
736  case CFT_Global:
737  case CFT_Device:
739  case CFT_HostDevice:
740  // An HD function counts as host code if we're compiling for host, and
741  // device code if we're compiling for device. Defer any errors in device
742  // mode until the function is known-emitted.
743  if (!getLangOpts().CUDAIsDevice)
745  if (IsLastErrorImmediate && Diags.getDiagnosticIDs()->isBuiltinNote(DiagID))
747  return (getEmissionStatus(CurFunContext) ==
751  default:
753  }
754  }();
755  return SemaDiagnosticBuilder(DiagKind, Loc, DiagID, CurFunContext, *this);
756 }
757 
759  unsigned DiagID) {
760  assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
761  FunctionDecl *CurFunContext = getCurFunctionDecl(/*AllowLambda=*/true);
762  SemaDiagnosticBuilder::Kind DiagKind = [&] {
763  if (!CurFunContext)
765  switch (CurrentCUDATarget()) {
766  case CFT_Host:
768  case CFT_HostDevice:
769  // An HD function counts as host code if we're compiling for host, and
770  // device code if we're compiling for device. Defer any errors in device
771  // mode until the function is known-emitted.
772  if (getLangOpts().CUDAIsDevice)
774  if (IsLastErrorImmediate && Diags.getDiagnosticIDs()->isBuiltinNote(DiagID))
776  return (getEmissionStatus(CurFunContext) ==
780  default:
782  }
783  }();
784  return SemaDiagnosticBuilder(DiagKind, Loc, DiagID, CurFunContext, *this);
785 }
786 
788  assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
789  assert(Callee && "Callee may not be null.");
790 
791  auto &ExprEvalCtx = ExprEvalContexts.back();
792  if (ExprEvalCtx.isUnevaluated() || ExprEvalCtx.isConstantEvaluated())
793  return true;
794 
795  // FIXME: Is bailing out early correct here? Should we instead assume that
796  // the caller is a global initializer?
797  FunctionDecl *Caller = getCurFunctionDecl(/*AllowLambda=*/true);
798  if (!Caller)
799  return true;
800 
801  // If the caller is known-emitted, mark the callee as known-emitted.
802  // Otherwise, mark the call in our call graph so we can traverse it later.
803  bool CallerKnownEmitted =
805  SemaDiagnosticBuilder::Kind DiagKind = [this, Caller, Callee,
806  CallerKnownEmitted] {
807  switch (IdentifyCUDAPreference(Caller, Callee)) {
808  case CFP_Never:
809  case CFP_WrongSide:
810  assert(Caller && "Never/wrongSide calls require a non-null caller");
811  // If we know the caller will be emitted, we know this wrong-side call
812  // will be emitted, so it's an immediate error. Otherwise, defer the
813  // error until we know the caller is emitted.
814  return CallerKnownEmitted
817  default:
819  }
820  }();
821 
822  if (DiagKind == SemaDiagnosticBuilder::K_Nop) {
823  // For -fgpu-rdc, keep track of external kernels used by host functions.
824  if (LangOpts.CUDAIsDevice && LangOpts.GPURelocatableDeviceCode &&
825  Callee->hasAttr<CUDAGlobalAttr>() && !Callee->isDefined())
827  return true;
828  }
829 
830  // Avoid emitting this error twice for the same location. Using a hashtable
831  // like this is unfortunate, but because we must continue parsing as normal
832  // after encountering a deferred error, it's otherwise very tricky for us to
833  // ensure that we only emit this deferred error once.
834  if (!LocsWithCUDACallDiags.insert({Caller, Loc}).second)
835  return true;
836 
837  SemaDiagnosticBuilder(DiagKind, Loc, diag::err_ref_bad_target, Caller, *this)
838  << IdentifyCUDATarget(Callee) << /*function*/ 0 << Callee
839  << IdentifyCUDATarget(Caller);
840  if (!Callee->getBuiltinID())
841  SemaDiagnosticBuilder(DiagKind, Callee->getLocation(),
842  diag::note_previous_decl, Caller, *this)
843  << Callee;
844  return DiagKind != SemaDiagnosticBuilder::K_Immediate &&
846 }
847 
848 // Check the wrong-sided reference capture of lambda for CUDA/HIP.
849 // A lambda function may capture a stack variable by reference when it is
850 // defined and uses the capture by reference when the lambda is called. When
851 // the capture and use happen on different sides, the capture is invalid and
852 // should be diagnosed.
854  const sema::Capture &Capture) {
855  // In host compilation we only need to check lambda functions emitted on host
856  // side. In such lambda functions, a reference capture is invalid only
857  // if the lambda structure is populated by a device function or kernel then
858  // is passed to and called by a host function. However that is impossible,
859  // since a device function or kernel can only call a device function, also a
860  // kernel cannot pass a lambda back to a host function since we cannot
861  // define a kernel argument type which can hold the lambda before the lambda
862  // itself is defined.
863  if (!LangOpts.CUDAIsDevice)
864  return;
865 
866  // File-scope lambda can only do init captures for global variables, which
867  // results in passing by value for these global variables.
868  FunctionDecl *Caller = getCurFunctionDecl(/*AllowLambda=*/true);
869  if (!Caller)
870  return;
871 
872  // In device compilation, we only need to check lambda functions which are
873  // emitted on device side. For such lambdas, a reference capture is invalid
874  // only if the lambda structure is populated by a host function then passed
875  // to and called in a device function or kernel.
876  bool CalleeIsDevice = Callee->hasAttr<CUDADeviceAttr>();
877  bool CallerIsHost =
878  !Caller->hasAttr<CUDAGlobalAttr>() && !Caller->hasAttr<CUDADeviceAttr>();
879  bool ShouldCheck = CalleeIsDevice && CallerIsHost;
880  if (!ShouldCheck || !Capture.isReferenceCapture())
881  return;
882  auto DiagKind = SemaDiagnosticBuilder::K_Deferred;
883  if (Capture.isVariableCapture()) {
885  diag::err_capture_bad_target, Callee, *this)
886  << Capture.getVariable();
887  } else if (Capture.isThisCapture()) {
888  // Capture of this pointer is allowed since this pointer may be pointing to
889  // managed memory which is accessible on both device and host sides. It only
890  // results in invalid memory access if this pointer points to memory not
891  // accessible on device side.
893  diag::warn_maybe_capture_bad_target_this_ptr, Callee,
894  *this);
895  }
896 }
897 
899  assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
900  if (Method->hasAttr<CUDAHostAttr>() || Method->hasAttr<CUDADeviceAttr>())
901  return;
902  Method->addAttr(CUDADeviceAttr::CreateImplicit(Context));
903  Method->addAttr(CUDAHostAttr::CreateImplicit(Context));
904 }
905 
907  const LookupResult &Previous) {
908  assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
909  CUDAFunctionTarget NewTarget = IdentifyCUDATarget(NewFD);
910  for (NamedDecl *OldND : Previous) {
911  FunctionDecl *OldFD = OldND->getAsFunction();
912  if (!OldFD)
913  continue;
914 
915  CUDAFunctionTarget OldTarget = IdentifyCUDATarget(OldFD);
916  // Don't allow HD and global functions to overload other functions with the
917  // same signature. We allow overloading based on CUDA attributes so that
918  // functions can have different implementations on the host and device, but
919  // HD/global functions "exist" in some sense on both the host and device, so
920  // should have the same implementation on both sides.
921  if (NewTarget != OldTarget &&
922  ((NewTarget == CFT_HostDevice) || (OldTarget == CFT_HostDevice) ||
923  (NewTarget == CFT_Global) || (OldTarget == CFT_Global)) &&
924  !IsOverload(NewFD, OldFD, /* UseMemberUsingDeclRules = */ false,
925  /* ConsiderCudaAttrs = */ false)) {
926  Diag(NewFD->getLocation(), diag::err_cuda_ovl_target)
927  << NewTarget << NewFD->getDeclName() << OldTarget << OldFD;
928  Diag(OldFD->getLocation(), diag::note_previous_declaration);
929  NewFD->setInvalidDecl();
930  break;
931  }
932  }
933 }
934 
935 template <typename AttrTy>
936 static void copyAttrIfPresent(Sema &S, FunctionDecl *FD,
937  const FunctionDecl &TemplateFD) {
938  if (AttrTy *Attribute = TemplateFD.getAttr<AttrTy>()) {
939  AttrTy *Clone = Attribute->clone(S.Context);
940  Clone->setInherited(true);
941  FD->addAttr(Clone);
942  }
943 }
944 
946  const FunctionTemplateDecl &TD) {
947  const FunctionDecl &TemplateFD = *TD.getTemplatedDecl();
948  copyAttrIfPresent<CUDAGlobalAttr>(*this, FD, TemplateFD);
949  copyAttrIfPresent<CUDAHostAttr>(*this, FD, TemplateFD);
950  copyAttrIfPresent<CUDADeviceAttr>(*this, FD, TemplateFD);
951 }
952 
954  if (getLangOpts().HIP)
955  return getLangOpts().HIPUseNewLaunchAPI ? "__hipPushCallConfiguration"
956  : "hipConfigureCall";
957 
958  // New CUDA kernel launch sequence.
961  return "__cudaPushCallConfiguration";
962 
963  // Legacy CUDA kernel configuration call
964  return "cudaConfigureCall";
965 }
clang::Language::CUDA
@ CUDA
clang::Sema::SpecialMemberOverloadResult
SpecialMemberOverloadResult - The overloading result for a special member function.
Definition: Sema.h:1379
clang::sema::Capture::isThisCapture
bool isThisCapture() const
Definition: ScopeInfo.h:611
clang::Sema::CFP_Never
@ CFP_Never
Definition: Sema.h:12519
clang::sema::Capture::getLocation
SourceLocation getLocation() const
Retrieve the location at which this variable was captured.
Definition: ScopeInfo.h:648
hasAttr
static bool hasAttr(const FunctionDecl *D, bool IgnoreImplicitAttr)
Definition: SemaCUDA.cpp:108
clang::CXXConstructorDecl
Represents a C++ constructor within a class.
Definition: DeclCXX.h:2421
clang::FunctionDecl::getNumParams
unsigned getNumParams() const
Return the number of parameters this function must have based on its FunctionType.
Definition: Decl.cpp:3442
clang::VarDecl::isFileVarDecl
bool isFileVarDecl() const
Returns true for file scoped variable declaration.
Definition: Decl.h:1257
clang::Sema::CFT_Host
@ CFT_Host
Definition: Sema.h:12486
clang::Sema::CUDAFunctionTarget
CUDAFunctionTarget
Definition: Sema.h:12483
string
string(SUBSTRING ${CMAKE_CURRENT_BINARY_DIR} 0 ${PATH_LIB_START} PATH_HEAD) string(SUBSTRING $
Definition: CMakeLists.txt:22
clang::Decl::hasAttr
bool hasAttr() const
Definition: DeclBase.h:542
clang::VarDecl::hasGlobalStorage
bool hasGlobalStorage() const
Returns true for all variables that do not have local storage.
Definition: Decl.h:1141
clang::Sema::getASTContext
ASTContext & getASTContext() const
Definition: Sema.h:1614
clang::QualType::isConstQualified
bool isConstQualified() const
Determine whether this type is const-qualified.
Definition: Type.h:6559
clang::Sema::SemaDiagnosticBuilder::Kind
Kind
Definition: Sema.h:1724
clang::Decl::getAsFunction
FunctionDecl * getAsFunction() LLVM_READONLY
Returns the function itself, or the templated function if this is a function template.
Definition: DeclBase.cpp:218
clang::FunctionDecl::isConstexpr
bool isConstexpr() const
Whether this is a (C++11) constexpr function or constexpr constructor.
Definition: Decl.h:2299
SemaInternal.h
clang::Sema::checkCUDATargetOverload
void checkCUDATargetOverload(FunctionDecl *NewFD, const LookupResult &Previous)
Check whether NewFD is a valid overload for CUDA.
Definition: SemaCUDA.cpp:906
hasImplicitAttr
static bool hasImplicitAttr(const FunctionDecl *D)
Definition: SemaCUDA.cpp:258
clang::sema::Capture::getVariable
VarDecl * getVariable() const
Definition: ScopeInfo.h:637
llvm::SmallVector
Definition: LLVM.h:38
clang::Sema::PushForceCUDAHostDevice
void PushForceCUDAHostDevice()
Increments our count of the number of times we've seen a pragma forcing functions to be host device.
Definition: SemaCUDA.cpp:37
Lookup.h
clang::Sema::CFP_WrongSide
@ CFP_WrongSide
Definition: Sema.h:12520
clang::SourceLocation
Encodes a location in the source.
Definition: SourceLocation.h:86
clang::sema::Capture::isReferenceCapture
bool isReferenceCapture() const
Definition: ScopeInfo.h:617
clang::ASTContext::CUDAConstantEvalContextRAII
Definition: ASTContext.h:670
clang::NamedDecl
This represents a decl that may have a name.
Definition: Decl.h:247
TargetInfo.h
clang::QualType
A (possibly-)qualified type.
Definition: Type.h:675
clang::FieldDecl
Represents a member of a struct/union/class.
Definition: Decl.h:2855
clang::LookupResult
Represents the results of name lookup.
Definition: Lookup.h:46
clang::ASTContext::getBaseElementType
QualType getBaseElementType(const ArrayType *VAT) const
Return the innermost element type of an array type.
Definition: ASTContext.cpp:6789
clang::TargetInfo::getSDKVersion
const llvm::VersionTuple & getSDKVersion() const
Definition: TargetInfo.h:1605
clang::Sema::Diag
SemaDiagnosticBuilder Diag(SourceLocation Loc, unsigned DiagID, bool DeferHint=false)
Emit a diagnostic.
Definition: Sema.cpp:1871
clang::Redeclarable::getFirstDecl
decl_type * getFirstDecl()
Return the first declaration of this declaration or itself if this is the only declaration.
Definition: Redeclarable.h:216
clang::Sema::CUDACheckLambdaCapture
void CUDACheckLambdaCapture(CXXMethodDecl *D, const sema::Capture &Capture)
Definition: SemaCUDA.cpp:853
clang::Sema::CVT_Unified
@ CVT_Unified
Emitted on both sides with different addresses.
Definition: Sema.h:12504
clang::UsingShadowDecl
Represents a shadow declaration implicitly introduced into a scope by a (resolved) using-declaration ...
Definition: DeclCXX.h:3207
clang::Sema::getSourceManager
SourceManager & getSourceManager() const
Definition: Sema.h:1612
llvm::Optional
Definition: LLVM.h:40
clang::Sema::isEmptyCudaDestructor
bool isEmptyCudaDestructor(SourceLocation Loc, CXXDestructorDecl *CD)
Definition: SemaCUDA.cpp:502
clang::FunctionDecl::isUserProvided
bool isUserProvided() const
True if this method is user-declared and was not deleted or defaulted on its first declaration.
Definition: Decl.h:2246
clang::Sema::BuildCallExpr
ExprResult BuildCallExpr(Scope *S, Expr *Fn, SourceLocation LParenLoc, MultiExprArg ArgExprs, SourceLocation RParenLoc, Expr *ExecConfig=nullptr, bool IsExecConfig=false, bool AllowRecovery=false)
BuildCallExpr - Handle a call to Fn with the specified array of arguments.
Definition: SemaExpr.cpp:6636
clang::Decl::getAttr
T * getAttr() const
Definition: DeclBase.h:538
clang::Sema::IsOverload
bool IsOverload(FunctionDecl *New, FunctionDecl *Old, bool IsForUsingDecl, bool ConsiderCudaAttrs=true, bool ConsiderRequiresClauses=true)
Definition: SemaOverload.cpp:1141
clang::sema::Capture
Definition: ScopeInfo.h:528
clang::DeclContext::getLexicalParent
DeclContext * getLexicalParent()
getLexicalParent - Returns the containing lexical DeclContext.
Definition: DeclBase.h:1892
clang::Sema::ExprEvalContexts
SmallVector< ExpressionEvaluationContextRecord, 8 > ExprEvalContexts
A stack of expression evaluation contexts.
Definition: Sema.h:1360
clang::ASTContext::getcudaConfigureCallDecl
FunctionDecl * getcudaConfigureCallDecl()
Definition: ASTContext.h:1406
clang::Sema::CVT_Device
@ CVT_Device
Definition: Sema.h:12501
clang::Sema::Context
ASTContext & Context
Definition: Sema.h:410
Preprocessor.h
clang::ExprError
ExprResult ExprError()
Definition: Ownership.h:278
Decl.h
clang::Sema::inferCUDATargetForImplicitSpecialMember
bool inferCUDATargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl, CXXSpecialMember CSM, CXXMethodDecl *MemberDecl, bool ConstRHS, bool Diagnose)
Given a implicit special member, infer its CUDA target from the calls it needs to make to underlying ...
Definition: SemaCUDA.cpp:324
clang::Sema::CUDADiagIfDeviceCode
SemaDiagnosticBuilder CUDADiagIfDeviceCode(SourceLocation Loc, unsigned DiagID)
Creates a SemaDiagnosticBuilder that emits the diagnostic if the current context is "used as device c...
Definition: SemaCUDA.cpp:728
clang::Sema::getLangOpts
const LangOptions & getLangOpts() const
Definition: Sema.h:1607
clang::Type::isCUDADeviceBuiltinTextureType
bool isCUDADeviceBuiltinTextureType() const
Check if the type is the CUDA device builtin texture type.
Definition: Type.cpp:4386
clang::Sema::getCurFunctionDecl
FunctionDecl * getCurFunctionDecl(bool AllowLambda=false)
Returns a pointer to the innermost enclosing function, or nullptr if the current context is not insid...
Definition: Sema.cpp:1451
clang::FunctionTemplateDecl
Declaration of a template function.
Definition: DeclTemplate.h:979
llvm::MutableArrayRef
Definition: LLVM.h:35
clang::Type::isReferenceType
bool isReferenceType() const
Definition: Type.h:6760
clang::CXXConstructorDecl::inits
init_range inits()
Definition: DeclCXX.h:2514
clang::VarDecl::hasInit
bool hasInit() const
Definition: Decl.cpp:2303
Template.h
clang::RecordType
A helper class that allows the use of isa/cast/dyncast to detect TagType objects of structs/unions/cl...
Definition: Type.h:4647
clang::Sema::MaybeAddCUDAConstantAttr
void MaybeAddCUDAConstantAttr(VarDecl *VD)
May add implicit CUDAConstantAttr attribute to VD, depending on VD and current compilation settings.
Definition: SemaCUDA.cpp:714
clang::FunctionDecl::isTrivial
bool isTrivial() const
Whether this function is "trivial" in some specialized C++ senses.
Definition: Decl.h:2223
clang::Sema::CFT_HostDevice
@ CFT_HostDevice
Definition: Sema.h:12487
clang::Sema::ActOnCUDAExecConfigExpr
ExprResult ActOnCUDAExecConfigExpr(Scope *S, SourceLocation LLLLoc, MultiExprArg ExecConfig, SourceLocation GGGLoc)
Definition: SemaCUDA.cpp:50
clang::Scope
Scope - A scope is a transient data structure that is used while parsing the program.
Definition: Scope.h:40
clang::Sema::SemaDiagnosticBuilder::K_Nop
@ K_Nop
Emit no diagnostics.
Definition: Sema.h:1726
clang::Sema::SemaDiagnosticBuilder::K_Immediate
@ K_Immediate
Emit the diagnostic immediately (i.e., behave like Sema::Diag()).
Definition: Sema.h:1728
clang::VarDecl::isStaticDataMember
bool isStaticDataMember() const
Determines whether this is a static data member.
Definition: Decl.h:1198
clang::FunctionDecl::hasTrivialBody
bool hasTrivialBody() const
Returns whether the function has a trivial body that does not require any specific codegen.
Definition: Decl.cpp:3023
resolveCalleeCUDATargetConflict
static bool resolveCalleeCUDATargetConflict(Sema::CUDAFunctionTarget Target1, Sema::CUDAFunctionTarget Target2, Sema::CUDAFunctionTarget *ResolvedTarget)
When an implicitly-declared special member has to invoke more than one base/field special member,...
Definition: SemaCUDA.cpp:304
clang::Sema::CFP_SameSide
@ CFP_SameSide
Definition: Sema.h:12524
clang::Sema::CFP_Native
@ CFP_Native
Definition: Sema.h:12526
clang::Type::getAs
const T * getAs() const
Member-template getAs<specific type>'.
Definition: Type.h:7243
clang::Decl::isInvalidDecl
bool isInvalidDecl() const
Definition: DeclBase.h:553
clang::Sema::Diags
DiagnosticsEngine & Diags
Definition: Sema.h:412
clang::CudaFeatureEnabled
bool CudaFeatureEnabled(llvm::VersionTuple, CudaFeature)
Definition: Cuda.cpp:261
clang::DiagnosticsEngine::getDiagnosticIDs
const IntrusiveRefCntPtr< DiagnosticIDs > & getDiagnosticIDs() const
Definition: Diagnostic.h:551
clang::CXXDestructorDecl
Represents a C++ destructor within a class.
Definition: DeclCXX.h:2688
clang::Sema::maybeAddCUDAHostDeviceAttrs
void maybeAddCUDAHostDeviceAttrs(FunctionDecl *FD, const LookupResult &Previous)
May add implicit CUDAHostAttr and CUDADeviceAttr attributes to FD, depending on FD and the current co...
Definition: SemaCUDA.cpp:662
clang::ParsedAttr
ParsedAttr - Represents a syntactic attribute.
Definition: ParsedAttr.h:234
clang::Sema::IsLastErrorImmediate
bool IsLastErrorImmediate
Is the last error level diagnostic immediate.
Definition: Sema.h:1826
clang::CudaFeature::CUDA_USES_NEW_LAUNCH
@ CUDA_USES_NEW_LAUNCH
ASTContext.h
clang::VarDecl
Represents a variable declaration or definition.
Definition: Decl.h:874
clang::Sema::isCUDAImplicitHostDeviceFunction
static bool isCUDAImplicitHostDeviceFunction(const FunctionDecl *D)
Definition: SemaCUDA.cpp:266
clang::Sema::SpecialMemberOverloadResult::getMethod
CXXMethodDecl * getMethod() const
Definition: Sema.h:1395
clang::Sema::CUDAVariableTarget
CUDAVariableTarget
Definition: Sema.h:12500
clang::Type::getAsCXXRecordDecl
CXXRecordDecl * getAsCXXRecordDecl() const
Retrieves the CXXRecordDecl that this type refers to, either because the type is a RecordType or beca...
Definition: Type.cpp:1759
clang::TagDecl::isUnion
bool isUnion() const
Definition: Decl.h:3554
clang::CXXRecordDecl::bases
base_class_range bases()
Definition: DeclCXX.h:596
ExprCXX.h
clang::Sema::CUDADiagIfHostCode
SemaDiagnosticBuilder CUDADiagIfHostCode(SourceLocation Loc, unsigned DiagID)
Creates a SemaDiagnosticBuilder that emits the diagnostic if the current context is "used as host cod...
Definition: SemaCUDA.cpp:758
clang::Sema::CFP_HostDevice
@ CFP_HostDevice
Definition: Sema.h:12523
clang::Sema::InstantiateFunctionDefinition
void InstantiateFunctionDefinition(SourceLocation PointOfInstantiation, FunctionDecl *Function, bool Recursive=false, bool DefinitionRequired=false, bool AtEndOfTU=false)
Instantiate the definition of the given function from its template.
Definition: SemaTemplateInstantiateDecl.cpp:4784
Sema.h
clang::CXXRecordDecl::vbases
base_class_range vbases()
Definition: DeclCXX.h:613
clang::Decl::isImplicit
bool isImplicit() const
isImplicit - Indicates whether the declaration was implicitly generated by the implementation.
Definition: DeclBase.h:558
clang::Sema::LangOpts
const LangOptions & LangOpts
Definition: Sema.h:408
clang::Sema::SemaDiagnosticBuilder::K_Deferred
@ K_Deferred
Create a deferred diagnostic, which is emitted only if the function it's attached to is codegen'ed.
Definition: Sema.h:1736
clang::VK_LValue
@ VK_LValue
An l-value expression is a reference to an object with independent storage.
Definition: Specifiers.h:124
clang::Sema::LocsWithCUDACallDiags
llvm::DenseSet< FunctionDeclAndLoc > LocsWithCUDACallDiags
FunctionDecls and SourceLocations for which CheckCUDACall has emitted a (maybe deferred) "bad call" d...
Definition: Sema.h:12400
clang::Sema::CVT_Both
@ CVT_Both
Emitted on host side only.
Definition: Sema.h:12503
clang::CXXRecordDecl
Represents a C++ struct/union/class.
Definition: DeclCXX.h:254
clang::Sema::PopForceCUDAHostDevice
bool PopForceCUDAHostDevice()
Decrements our count of the number of times we've seen a pragma forcing functions to be host device.
Definition: SemaCUDA.cpp:42
clang::Sema::CUDAFunctionPreference
CUDAFunctionPreference
Definition: Sema.h:12518
clang::Type::isDependentType
bool isDependentType() const
Whether this type is a dependent type, meaning that its definition somehow depends on a template para...
Definition: Type.h:2185
clang::Decl::hasAttrs
bool hasAttrs() const
Definition: DeclBase.h:484
clang::Sema::CUDASetLambdaAttrs
void CUDASetLambdaAttrs(CXXMethodDecl *Method)
Set device or host device attributes on the given lambda operator() method.
Definition: SemaCUDA.cpp:898
SemaDiagnostic.h
clang::RecordDecl::fields
field_range fields() const
Definition: Decl.h:4127
hasExplicitAttr
static bool hasExplicitAttr(const VarDecl *D)
Definition: SemaCUDA.cpp:29
clang::Sema::isEmptyCudaConstructor
bool isEmptyCudaConstructor(SourceLocation Loc, CXXConstructorDecl *CD)
Definition: SemaCUDA.cpp:464
clang::Sema::SemaDiagnosticBuilder::K_ImmediateWithCallStack
@ K_ImmediateWithCallStack
Emit the diagnostic immediately, and, if it's a warning or error, also emit a call stack showing how ...
Definition: Sema.h:1732
clang::Sema
Sema - This implements semantic analysis and AST building for C.
Definition: Sema.h:354
copyAttrIfPresent
static void copyAttrIfPresent(Sema &S, FunctionDecl *FD, const FunctionDecl &TemplateFD)
Definition: SemaCUDA.cpp:936
clang::Sema::CFT_Device
@ CFT_Device
Definition: Sema.h:12484
clang::Sema::ContextRAII
A RAII object to temporarily push a declaration context.
Definition: Sema.h:985
clang::VarDecl::getInit
const Expr * getInit() const
Definition: Decl.h:1283
clang::ASTContext::getTargetInfo
const TargetInfo & getTargetInfo() const
Definition: ASTContext.h:768
clang::ASTContext::CUDAExternalDeviceDeclODRUsedByHost
llvm::DenseSet< const ValueDecl * > CUDAExternalDeviceDeclODRUsedByHost
Keep track of CUDA/HIP external kernels or device variables ODR-used by host code.
Definition: ASTContext.h:1165
clang::Sema::SemaDiagnosticBuilder
A generic diagnostic builder for errors which may or may not be deferred.
Definition: Sema.h:1722
clang::ActionResult< Expr * >
clang::FunctionTemplateDecl::getTemplatedDecl
FunctionDecl * getTemplatedDecl() const
Get the underlying function declaration of the template.
Definition: DeclTemplate.h:1034
clang::Sema::getEmissionStatus
FunctionEmissionStatus getEmissionStatus(FunctionDecl *Decl, bool Final=false)
Definition: SemaDecl.cpp:19039
ScopeInfo.h
Cuda.h
clang::Sema::FunctionEmissionStatus::Emitted
@ Emitted
clang
Definition: CalledOnceCheck.h:17
clang::Sema::CFT_Global
@ CFT_Global
Definition: Sema.h:12485
clang::VarDecl::isConstexpr
bool isConstexpr() const
Whether this variable is (C++11) constexpr.
Definition: Decl.h:1477
clang::Sema::IdentifyCUDAPreference
CUDAFunctionPreference IdentifyCUDAPreference(const FunctionDecl *Caller, const FunctionDecl *Callee)
Identifies relative preference of a given Caller/Callee combination, based on their host/device attri...
Definition: SemaCUDA.cpp:207
clang::Sema::EraseUnwantedCUDAMatches
void EraseUnwantedCUDAMatches(const FunctionDecl *Caller, SmallVectorImpl< std::pair< DeclAccessPair, FunctionDecl * >> &Matches)
Finds a function in Matches with highest calling priority from Caller context and erases all function...
Definition: SemaCUDA.cpp:272
clang::CXXBaseSpecifier
Represents a base class of a C++ class.
Definition: DeclCXX.h:146
clang::Attr
Attr - This represents one attribute.
Definition: Attr.h:41
clang::NamedDecl::getDeclName
DeclarationName getDeclName() const
Get the actual, stored name of the declaration, which may be a special name.
Definition: Decl.h:311
clang::CXXRecordDecl::isAbstract
bool isAbstract() const
Determine whether this class has a pure virtual function.
Definition: DeclCXX.h:1176
unsigned
clang::FunctionDecl::isVariadic
bool isVariadic() const
Whether this function is variadic.
Definition: Decl.cpp:2980
clang::Decl::getAttrs
AttrVec & getAttrs()
Definition: DeclBase.h:490
clang::Sema::CurrentCUDATarget
CUDAFunctionTarget CurrentCUDATarget()
Gets the CUDA target for the current context.
Definition: Sema.h:12510
clang::Sema::getCudaConfigureFuncName
std::string getCudaConfigureFuncName() const
Returns the name of the launch configuration function.
Definition: SemaCUDA.cpp:953
clang::CXXRecordDecl::isDynamicClass
bool isDynamicClass() const
Definition: DeclCXX.h:562
clang::RecordType::getDecl
RecordDecl * getDecl() const
Definition: Type.h:4657
clang::Sema::LookupSpecialMember
SpecialMemberOverloadResult LookupSpecialMember(CXXRecordDecl *D, CXXSpecialMember SM, bool ConstArg, bool VolatileArg, bool RValueThis, bool ConstThis, bool VolatileThis)
Definition: SemaLookup.cpp:3085
clang::Sema::inheritCUDATargetAttrs
void inheritCUDATargetAttrs(FunctionDecl *FD, const FunctionTemplateDecl &TD)
Copies target attributes from the template TD to the function FD.
Definition: SemaCUDA.cpp:945
clang::Sema::CXXSpecialMember
CXXSpecialMember
Kinds of C++ special members.
Definition: Sema.h:1502
clang::Decl::setInvalidDecl
void setInvalidDecl(bool Invalid=true)
setInvalidDecl - Indicates the Decl had a semantic error.
Definition: DeclBase.cpp:132
clang::Language::HIP
@ HIP
llvm::SmallVectorImpl
Definition: Randstruct.h:18
clang::FunctionDecl::isTemplateInstantiation
bool isTemplateInstantiation() const
Determines if the given function was instantiated from a function template.
Definition: Decl.cpp:3818
clang::ValueDecl::getType
QualType getType() const
Definition: Decl.h:685
clang::Sema::CVT_Host
@ CVT_Host
Emitted on device side with a shadow variable on host side.
Definition: Sema.h:12502
Previous
StateNode * Previous
Definition: UnwrappedLineFormatter.cpp:1092
clang::Expr
This represents one expression.
Definition: Expr.h:109
clang::ParsedAttributesView
Definition: ParsedAttr.h:898
clang::Sema::CFT_InvalidTarget
@ CFT_InvalidTarget
Definition: Sema.h:12488
clang::sema::Capture::isVariableCapture
bool isVariableCapture() const
Definition: ScopeInfo.h:612
clang::CXXCtorInitializer
Represents a C++ base or member initializer.
Definition: DeclCXX.h:2192
clang::Sema::checkAllowedCUDAInitializer
void checkAllowedCUDAInitializer(VarDecl *VD)
Definition: SemaCUDA.cpp:610
clang::Sema::MarkFunctionReferenced
void MarkFunctionReferenced(SourceLocation Loc, FunctionDecl *Func, bool MightBeOdrUse=true)
Mark a function referenced, and check whether it is odr-used (C++ [basic.def.odr]p2,...
Definition: SemaExpr.cpp:17799
clang::Decl::getLocation
SourceLocation getLocation() const
Definition: DeclBase.h:425
clang::Decl::addAttr
void addAttr(Attr *A)
Definition: DeclBase.cpp:886
clang::DeclRefExpr
A reference to a declared variable, function, enum, etc.
Definition: Expr.h:1223
clang::FunctionDecl
Represents a function declaration or definition.
Definition: Decl.h:1872
clang::Type::isCUDADeviceBuiltinSurfaceType
bool isCUDADeviceBuiltinSurfaceType() const
Check if the type is the CUDA device builtin surface type.
Definition: Type.cpp:4379
clang::CallExpr
CallExpr - Represents a function call (C99 6.5.2.2, C++ [expr.call]).
Definition: Expr.h:2801
clang::CXXConstructExpr
Represents a call to a C++ constructor.
Definition: ExprCXX.h:1460
clang::FunctionDecl::isDefined
bool isDefined(const FunctionDecl *&Definition, bool CheckForPendingFriendDefinition=false) const
Returns true if the function has a definition that does not need to be instantiated.
Definition: Decl.cpp:3059
clang::Sema::IdentifyCUDATarget
CUDAFunctionTarget IdentifyCUDATarget(const FunctionDecl *D, bool IgnoreImplicitHDAttr=false)
Determines whether the given function is a CUDA device/host/kernel/etc.
Definition: SemaCUDA.cpp:116
clang::Sema::CheckCUDACall
bool CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee)
Check whether we're allowed to call Callee from the current context.
Definition: SemaCUDA.cpp:787
clang::CXXMethodDecl::getParent
const CXXRecordDecl * getParent() const
Return the parent of this method declaration, which is the class in which this method is defined.
Definition: DeclCXX.h:2083
clang::CXXMethodDecl
Represents a static or instance method of a struct/union/class.
Definition: DeclCXX.h:1968
clang::Decl::getDeclContext
DeclContext * getDeclContext()
Definition: DeclBase.h:434