xref: /freebsd/contrib/llvm-project/clang/lib/Sema/SemaCUDA.cpp (revision 5b56413d04e608379c9a306373554a8e4d321bc0)
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"
22 #include "clang/Sema/SemaDiagnostic.h"
23 #include "clang/Sema/SemaInternal.h"
24 #include "clang/Sema/Template.h"
25 #include "llvm/ADT/SmallVector.h"
26 #include <optional>
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 
37 void Sema::PushForceCUDAHostDevice() {
38   assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
39   ForceCUDAHostDeviceDepth++;
40 }
41 
42 bool Sema::PopForceCUDAHostDevice() {
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 
50 ExprResult Sema::ActOnCUDAExecConfigExpr(Scope *S, SourceLocation LLLLoc,
51                                          MultiExprArg ExecConfig,
52                                          SourceLocation GGGLoc) {
53   FunctionDecl *ConfigDecl = Context.getcudaConfigureCallDecl();
54   if (!ConfigDecl)
55     return ExprError(Diag(LLLLoc, diag::err_undeclared_var_use)
56                      << getCudaConfigureFuncName());
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 
67 Sema::CUDAFunctionTarget
68 Sema::IdentifyCUDATarget(const ParsedAttributesView &Attrs) {
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 Decl *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 Sema::CUDATargetContextRAII::CUDATargetContextRAII(Sema &S_,
116                                                    CUDATargetContextKind K,
117                                                    Decl *D)
118     : S(S_) {
119   SavedCtx = S.CurCUDATargetCtx;
120   assert(K == CTCK_InitGlobalVar);
121   auto *VD = dyn_cast_or_null<VarDecl>(D);
122   if (VD && VD->hasGlobalStorage() && !VD->isStaticLocal()) {
123     auto Target = CFT_Host;
124     if ((hasAttr<CUDADeviceAttr>(VD, /*IgnoreImplicit=*/true) &&
125          !hasAttr<CUDAHostAttr>(VD, /*IgnoreImplicit=*/true)) ||
126         hasAttr<CUDASharedAttr>(VD, /*IgnoreImplicit=*/true) ||
127         hasAttr<CUDAConstantAttr>(VD, /*IgnoreImplicit=*/true))
128       Target = CFT_Device;
129     S.CurCUDATargetCtx = {Target, K, VD};
130   }
131 }
132 
133 /// IdentifyCUDATarget - Determine the CUDA compilation target for this function
134 Sema::CUDAFunctionTarget Sema::IdentifyCUDATarget(const FunctionDecl *D,
135                                                   bool IgnoreImplicitHDAttr) {
136   // Code that lives outside a function gets the target from CurCUDATargetCtx.
137   if (D == nullptr)
138     return CurCUDATargetCtx.Target;
139 
140   if (D->hasAttr<CUDAInvalidTargetAttr>())
141     return CFT_InvalidTarget;
142 
143   if (D->hasAttr<CUDAGlobalAttr>())
144     return CFT_Global;
145 
146   if (hasAttr<CUDADeviceAttr>(D, IgnoreImplicitHDAttr)) {
147     if (hasAttr<CUDAHostAttr>(D, IgnoreImplicitHDAttr))
148       return CFT_HostDevice;
149     return CFT_Device;
150   } else if (hasAttr<CUDAHostAttr>(D, IgnoreImplicitHDAttr)) {
151     return CFT_Host;
152   } else if ((D->isImplicit() || !D->isUserProvided()) &&
153              !IgnoreImplicitHDAttr) {
154     // Some implicit declarations (like intrinsic functions) are not marked.
155     // Set the most lenient target on them for maximal flexibility.
156     return CFT_HostDevice;
157   }
158 
159   return CFT_Host;
160 }
161 
162 /// IdentifyTarget - Determine the CUDA compilation target for this variable.
163 Sema::CUDAVariableTarget Sema::IdentifyCUDATarget(const VarDecl *Var) {
164   if (Var->hasAttr<HIPManagedAttr>())
165     return CVT_Unified;
166   // Only constexpr and const variabless with implicit constant attribute
167   // are emitted on both sides. Such variables are promoted to device side
168   // only if they have static constant intializers on device side.
169   if ((Var->isConstexpr() || Var->getType().isConstQualified()) &&
170       Var->hasAttr<CUDAConstantAttr>() &&
171       !hasExplicitAttr<CUDAConstantAttr>(Var))
172     return CVT_Both;
173   if (Var->hasAttr<CUDADeviceAttr>() || Var->hasAttr<CUDAConstantAttr>() ||
174       Var->hasAttr<CUDASharedAttr>() ||
175       Var->getType()->isCUDADeviceBuiltinSurfaceType() ||
176       Var->getType()->isCUDADeviceBuiltinTextureType())
177     return CVT_Device;
178   // Function-scope static variable without explicit device or constant
179   // attribute are emitted
180   //  - on both sides in host device functions
181   //  - on device side in device or global functions
182   if (auto *FD = dyn_cast<FunctionDecl>(Var->getDeclContext())) {
183     switch (IdentifyCUDATarget(FD)) {
184     case CFT_HostDevice:
185       return CVT_Both;
186     case CFT_Device:
187     case CFT_Global:
188       return CVT_Device;
189     default:
190       return CVT_Host;
191     }
192   }
193   return CVT_Host;
194 }
195 
196 // * CUDA Call preference table
197 //
198 // F - from,
199 // T - to
200 // Ph - preference in host mode
201 // Pd - preference in device mode
202 // H  - handled in (x)
203 // Preferences: N:native, SS:same side, HD:host-device, WS:wrong side, --:never.
204 //
205 // | F  | T  | Ph  | Pd  |  H  |
206 // |----+----+-----+-----+-----+
207 // | d  | d  | N   | N   | (c) |
208 // | d  | g  | --  | --  | (a) |
209 // | d  | h  | --  | --  | (e) |
210 // | d  | hd | HD  | HD  | (b) |
211 // | g  | d  | N   | N   | (c) |
212 // | g  | g  | --  | --  | (a) |
213 // | g  | h  | --  | --  | (e) |
214 // | g  | hd | HD  | HD  | (b) |
215 // | h  | d  | --  | --  | (e) |
216 // | h  | g  | N   | N   | (c) |
217 // | h  | h  | N   | N   | (c) |
218 // | h  | hd | HD  | HD  | (b) |
219 // | hd | d  | WS  | SS  | (d) |
220 // | hd | g  | SS  | --  |(d/a)|
221 // | hd | h  | SS  | WS  | (d) |
222 // | hd | hd | HD  | HD  | (b) |
223 
224 Sema::CUDAFunctionPreference
225 Sema::IdentifyCUDAPreference(const FunctionDecl *Caller,
226                              const FunctionDecl *Callee) {
227   assert(Callee && "Callee must be valid.");
228 
229   // Treat ctor/dtor as host device function in device var initializer to allow
230   // trivial ctor/dtor without device attr to be used. Non-trivial ctor/dtor
231   // will be diagnosed by checkAllowedCUDAInitializer.
232   if (Caller == nullptr && CurCUDATargetCtx.Kind == CTCK_InitGlobalVar &&
233       CurCUDATargetCtx.Target == CFT_Device &&
234       (isa<CXXConstructorDecl>(Callee) || isa<CXXDestructorDecl>(Callee)))
235     return CFP_HostDevice;
236 
237   CUDAFunctionTarget CallerTarget = IdentifyCUDATarget(Caller);
238   CUDAFunctionTarget CalleeTarget = IdentifyCUDATarget(Callee);
239 
240   // If one of the targets is invalid, the check always fails, no matter what
241   // the other target is.
242   if (CallerTarget == CFT_InvalidTarget || CalleeTarget == CFT_InvalidTarget)
243     return CFP_Never;
244 
245   // (a) Can't call global from some contexts until we support CUDA's
246   // dynamic parallelism.
247   if (CalleeTarget == CFT_Global &&
248       (CallerTarget == CFT_Global || CallerTarget == CFT_Device))
249     return CFP_Never;
250 
251   // (b) Calling HostDevice is OK for everyone.
252   if (CalleeTarget == CFT_HostDevice)
253     return CFP_HostDevice;
254 
255   // (c) Best case scenarios
256   if (CalleeTarget == CallerTarget ||
257       (CallerTarget == CFT_Host && CalleeTarget == CFT_Global) ||
258       (CallerTarget == CFT_Global && CalleeTarget == CFT_Device))
259     return CFP_Native;
260 
261   // HipStdPar mode is special, in that assessing whether a device side call to
262   // a host target is deferred to a subsequent pass, and cannot unambiguously be
263   // adjudicated in the AST, hence we optimistically allow them to pass here.
264   if (getLangOpts().HIPStdPar &&
265       (CallerTarget == CFT_Global || CallerTarget == CFT_Device ||
266        CallerTarget == CFT_HostDevice) &&
267       CalleeTarget == CFT_Host)
268     return CFP_HostDevice;
269 
270   // (d) HostDevice behavior depends on compilation mode.
271   if (CallerTarget == CFT_HostDevice) {
272     // It's OK to call a compilation-mode matching function from an HD one.
273     if ((getLangOpts().CUDAIsDevice && CalleeTarget == CFT_Device) ||
274         (!getLangOpts().CUDAIsDevice &&
275          (CalleeTarget == CFT_Host || CalleeTarget == CFT_Global)))
276       return CFP_SameSide;
277 
278     // Calls from HD to non-mode-matching functions (i.e., to host functions
279     // when compiling in device mode or to device functions when compiling in
280     // host mode) are allowed at the sema level, but eventually rejected if
281     // they're ever codegened.  TODO: Reject said calls earlier.
282     return CFP_WrongSide;
283   }
284 
285   // (e) Calling across device/host boundary is not something you should do.
286   if ((CallerTarget == CFT_Host && CalleeTarget == CFT_Device) ||
287       (CallerTarget == CFT_Device && CalleeTarget == CFT_Host) ||
288       (CallerTarget == CFT_Global && CalleeTarget == CFT_Host))
289     return CFP_Never;
290 
291   llvm_unreachable("All cases should've been handled by now.");
292 }
293 
294 template <typename AttrT> static bool hasImplicitAttr(const FunctionDecl *D) {
295   if (!D)
296     return false;
297   if (auto *A = D->getAttr<AttrT>())
298     return A->isImplicit();
299   return D->isImplicit();
300 }
301 
302 bool Sema::isCUDAImplicitHostDeviceFunction(const FunctionDecl *D) {
303   bool IsImplicitDevAttr = hasImplicitAttr<CUDADeviceAttr>(D);
304   bool IsImplicitHostAttr = hasImplicitAttr<CUDAHostAttr>(D);
305   return IsImplicitDevAttr && IsImplicitHostAttr;
306 }
307 
308 void Sema::EraseUnwantedCUDAMatches(
309     const FunctionDecl *Caller,
310     SmallVectorImpl<std::pair<DeclAccessPair, FunctionDecl *>> &Matches) {
311   if (Matches.size() <= 1)
312     return;
313 
314   using Pair = std::pair<DeclAccessPair, FunctionDecl*>;
315 
316   // Gets the CUDA function preference for a call from Caller to Match.
317   auto GetCFP = [&](const Pair &Match) {
318     return IdentifyCUDAPreference(Caller, Match.second);
319   };
320 
321   // Find the best call preference among the functions in Matches.
322   CUDAFunctionPreference BestCFP = GetCFP(*std::max_element(
323       Matches.begin(), Matches.end(),
324       [&](const Pair &M1, const Pair &M2) { return GetCFP(M1) < GetCFP(M2); }));
325 
326   // Erase all functions with lower priority.
327   llvm::erase_if(Matches,
328                  [&](const Pair &Match) { return GetCFP(Match) < BestCFP; });
329 }
330 
331 /// When an implicitly-declared special member has to invoke more than one
332 /// base/field special member, conflicts may occur in the targets of these
333 /// members. For example, if one base's member __host__ and another's is
334 /// __device__, it's a conflict.
335 /// This function figures out if the given targets \param Target1 and
336 /// \param Target2 conflict, and if they do not it fills in
337 /// \param ResolvedTarget with a target that resolves for both calls.
338 /// \return true if there's a conflict, false otherwise.
339 static bool
340 resolveCalleeCUDATargetConflict(Sema::CUDAFunctionTarget Target1,
341                                 Sema::CUDAFunctionTarget Target2,
342                                 Sema::CUDAFunctionTarget *ResolvedTarget) {
343   // Only free functions and static member functions may be global.
344   assert(Target1 != Sema::CFT_Global);
345   assert(Target2 != Sema::CFT_Global);
346 
347   if (Target1 == Sema::CFT_HostDevice) {
348     *ResolvedTarget = Target2;
349   } else if (Target2 == Sema::CFT_HostDevice) {
350     *ResolvedTarget = Target1;
351   } else if (Target1 != Target2) {
352     return true;
353   } else {
354     *ResolvedTarget = Target1;
355   }
356 
357   return false;
358 }
359 
360 bool Sema::inferCUDATargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl,
361                                                    CXXSpecialMember CSM,
362                                                    CXXMethodDecl *MemberDecl,
363                                                    bool ConstRHS,
364                                                    bool Diagnose) {
365   // If the defaulted special member is defined lexically outside of its
366   // owning class, or the special member already has explicit device or host
367   // attributes, do not infer.
368   bool InClass = MemberDecl->getLexicalParent() == MemberDecl->getParent();
369   bool HasH = MemberDecl->hasAttr<CUDAHostAttr>();
370   bool HasD = MemberDecl->hasAttr<CUDADeviceAttr>();
371   bool HasExplicitAttr =
372       (HasD && !MemberDecl->getAttr<CUDADeviceAttr>()->isImplicit()) ||
373       (HasH && !MemberDecl->getAttr<CUDAHostAttr>()->isImplicit());
374   if (!InClass || HasExplicitAttr)
375     return false;
376 
377   std::optional<CUDAFunctionTarget> InferredTarget;
378 
379   // We're going to invoke special member lookup; mark that these special
380   // members are called from this one, and not from its caller.
381   ContextRAII MethodContext(*this, MemberDecl);
382 
383   // Look for special members in base classes that should be invoked from here.
384   // Infer the target of this member base on the ones it should call.
385   // Skip direct and indirect virtual bases for abstract classes.
386   llvm::SmallVector<const CXXBaseSpecifier *, 16> Bases;
387   for (const auto &B : ClassDecl->bases()) {
388     if (!B.isVirtual()) {
389       Bases.push_back(&B);
390     }
391   }
392 
393   if (!ClassDecl->isAbstract()) {
394     llvm::append_range(Bases, llvm::make_pointer_range(ClassDecl->vbases()));
395   }
396 
397   for (const auto *B : Bases) {
398     const RecordType *BaseType = B->getType()->getAs<RecordType>();
399     if (!BaseType) {
400       continue;
401     }
402 
403     CXXRecordDecl *BaseClassDecl = cast<CXXRecordDecl>(BaseType->getDecl());
404     Sema::SpecialMemberOverloadResult SMOR =
405         LookupSpecialMember(BaseClassDecl, CSM,
406                             /* ConstArg */ ConstRHS,
407                             /* VolatileArg */ false,
408                             /* RValueThis */ false,
409                             /* ConstThis */ false,
410                             /* VolatileThis */ false);
411 
412     if (!SMOR.getMethod())
413       continue;
414 
415     CUDAFunctionTarget BaseMethodTarget = IdentifyCUDATarget(SMOR.getMethod());
416     if (!InferredTarget) {
417       InferredTarget = BaseMethodTarget;
418     } else {
419       bool ResolutionError = resolveCalleeCUDATargetConflict(
420           *InferredTarget, BaseMethodTarget, &*InferredTarget);
421       if (ResolutionError) {
422         if (Diagnose) {
423           Diag(ClassDecl->getLocation(),
424                diag::note_implicit_member_target_infer_collision)
425               << (unsigned)CSM << *InferredTarget << BaseMethodTarget;
426         }
427         MemberDecl->addAttr(CUDAInvalidTargetAttr::CreateImplicit(Context));
428         return true;
429       }
430     }
431   }
432 
433   // Same as for bases, but now for special members of fields.
434   for (const auto *F : ClassDecl->fields()) {
435     if (F->isInvalidDecl()) {
436       continue;
437     }
438 
439     const RecordType *FieldType =
440         Context.getBaseElementType(F->getType())->getAs<RecordType>();
441     if (!FieldType) {
442       continue;
443     }
444 
445     CXXRecordDecl *FieldRecDecl = cast<CXXRecordDecl>(FieldType->getDecl());
446     Sema::SpecialMemberOverloadResult SMOR =
447         LookupSpecialMember(FieldRecDecl, CSM,
448                             /* ConstArg */ ConstRHS && !F->isMutable(),
449                             /* VolatileArg */ false,
450                             /* RValueThis */ false,
451                             /* ConstThis */ false,
452                             /* VolatileThis */ false);
453 
454     if (!SMOR.getMethod())
455       continue;
456 
457     CUDAFunctionTarget FieldMethodTarget =
458         IdentifyCUDATarget(SMOR.getMethod());
459     if (!InferredTarget) {
460       InferredTarget = FieldMethodTarget;
461     } else {
462       bool ResolutionError = resolveCalleeCUDATargetConflict(
463           *InferredTarget, FieldMethodTarget, &*InferredTarget);
464       if (ResolutionError) {
465         if (Diagnose) {
466           Diag(ClassDecl->getLocation(),
467                diag::note_implicit_member_target_infer_collision)
468               << (unsigned)CSM << *InferredTarget << FieldMethodTarget;
469         }
470         MemberDecl->addAttr(CUDAInvalidTargetAttr::CreateImplicit(Context));
471         return true;
472       }
473     }
474   }
475 
476 
477   // If no target was inferred, mark this member as __host__ __device__;
478   // it's the least restrictive option that can be invoked from any target.
479   bool NeedsH = true, NeedsD = true;
480   if (InferredTarget) {
481     if (*InferredTarget == CFT_Device)
482       NeedsH = false;
483     else if (*InferredTarget == CFT_Host)
484       NeedsD = false;
485   }
486 
487   // We either setting attributes first time, or the inferred ones must match
488   // previously set ones.
489   if (NeedsD && !HasD)
490     MemberDecl->addAttr(CUDADeviceAttr::CreateImplicit(Context));
491   if (NeedsH && !HasH)
492     MemberDecl->addAttr(CUDAHostAttr::CreateImplicit(Context));
493 
494   return false;
495 }
496 
497 bool Sema::isEmptyCudaConstructor(SourceLocation Loc, CXXConstructorDecl *CD) {
498   if (!CD->isDefined() && CD->isTemplateInstantiation())
499     InstantiateFunctionDefinition(Loc, CD->getFirstDecl());
500 
501   // (E.2.3.1, CUDA 7.5) A constructor for a class type is considered
502   // empty at a point in the translation unit, if it is either a
503   // trivial constructor
504   if (CD->isTrivial())
505     return true;
506 
507   // ... or it satisfies all of the following conditions:
508   // The constructor function has been defined.
509   // The constructor function has no parameters,
510   // and the function body is an empty compound statement.
511   if (!(CD->hasTrivialBody() && CD->getNumParams() == 0))
512     return false;
513 
514   // Its class has no virtual functions and no virtual base classes.
515   if (CD->getParent()->isDynamicClass())
516     return false;
517 
518   // Union ctor does not call ctors of its data members.
519   if (CD->getParent()->isUnion())
520     return true;
521 
522   // The only form of initializer allowed is an empty constructor.
523   // This will recursively check all base classes and member initializers
524   if (!llvm::all_of(CD->inits(), [&](const CXXCtorInitializer *CI) {
525         if (const CXXConstructExpr *CE =
526                 dyn_cast<CXXConstructExpr>(CI->getInit()))
527           return isEmptyCudaConstructor(Loc, CE->getConstructor());
528         return false;
529       }))
530     return false;
531 
532   return true;
533 }
534 
535 bool Sema::isEmptyCudaDestructor(SourceLocation Loc, CXXDestructorDecl *DD) {
536   // No destructor -> no problem.
537   if (!DD)
538     return true;
539 
540   if (!DD->isDefined() && DD->isTemplateInstantiation())
541     InstantiateFunctionDefinition(Loc, DD->getFirstDecl());
542 
543   // (E.2.3.1, CUDA 7.5) A destructor for a class type is considered
544   // empty at a point in the translation unit, if it is either a
545   // trivial constructor
546   if (DD->isTrivial())
547     return true;
548 
549   // ... or it satisfies all of the following conditions:
550   // The destructor function has been defined.
551   // and the function body is an empty compound statement.
552   if (!DD->hasTrivialBody())
553     return false;
554 
555   const CXXRecordDecl *ClassDecl = DD->getParent();
556 
557   // Its class has no virtual functions and no virtual base classes.
558   if (ClassDecl->isDynamicClass())
559     return false;
560 
561   // Union does not have base class and union dtor does not call dtors of its
562   // data members.
563   if (DD->getParent()->isUnion())
564     return true;
565 
566   // Only empty destructors are allowed. This will recursively check
567   // destructors for all base classes...
568   if (!llvm::all_of(ClassDecl->bases(), [&](const CXXBaseSpecifier &BS) {
569         if (CXXRecordDecl *RD = BS.getType()->getAsCXXRecordDecl())
570           return isEmptyCudaDestructor(Loc, RD->getDestructor());
571         return true;
572       }))
573     return false;
574 
575   // ... and member fields.
576   if (!llvm::all_of(ClassDecl->fields(), [&](const FieldDecl *Field) {
577         if (CXXRecordDecl *RD = Field->getType()
578                                     ->getBaseElementTypeUnsafe()
579                                     ->getAsCXXRecordDecl())
580           return isEmptyCudaDestructor(Loc, RD->getDestructor());
581         return true;
582       }))
583     return false;
584 
585   return true;
586 }
587 
588 namespace {
589 enum CUDAInitializerCheckKind {
590   CICK_DeviceOrConstant, // Check initializer for device/constant variable
591   CICK_Shared,           // Check initializer for shared variable
592 };
593 
594 bool IsDependentVar(VarDecl *VD) {
595   if (VD->getType()->isDependentType())
596     return true;
597   if (const auto *Init = VD->getInit())
598     return Init->isValueDependent();
599   return false;
600 }
601 
602 // Check whether a variable has an allowed initializer for a CUDA device side
603 // variable with global storage. \p VD may be a host variable to be checked for
604 // potential promotion to device side variable.
605 //
606 // CUDA/HIP allows only empty constructors as initializers for global
607 // variables (see E.2.3.1, CUDA 7.5). The same restriction also applies to all
608 // __shared__ variables whether they are local or not (they all are implicitly
609 // static in CUDA). One exception is that CUDA allows constant initializers
610 // for __constant__ and __device__ variables.
611 bool HasAllowedCUDADeviceStaticInitializer(Sema &S, VarDecl *VD,
612                                            CUDAInitializerCheckKind CheckKind) {
613   assert(!VD->isInvalidDecl() && VD->hasGlobalStorage());
614   assert(!IsDependentVar(VD) && "do not check dependent var");
615   const Expr *Init = VD->getInit();
616   auto IsEmptyInit = [&](const Expr *Init) {
617     if (!Init)
618       return true;
619     if (const auto *CE = dyn_cast<CXXConstructExpr>(Init)) {
620       return S.isEmptyCudaConstructor(VD->getLocation(), CE->getConstructor());
621     }
622     return false;
623   };
624   auto IsConstantInit = [&](const Expr *Init) {
625     assert(Init);
626     ASTContext::CUDAConstantEvalContextRAII EvalCtx(S.Context,
627                                                     /*NoWronSidedVars=*/true);
628     return Init->isConstantInitializer(S.Context,
629                                        VD->getType()->isReferenceType());
630   };
631   auto HasEmptyDtor = [&](VarDecl *VD) {
632     if (const auto *RD = VD->getType()->getAsCXXRecordDecl())
633       return S.isEmptyCudaDestructor(VD->getLocation(), RD->getDestructor());
634     return true;
635   };
636   if (CheckKind == CICK_Shared)
637     return IsEmptyInit(Init) && HasEmptyDtor(VD);
638   return S.LangOpts.GPUAllowDeviceInit ||
639          ((IsEmptyInit(Init) || IsConstantInit(Init)) && HasEmptyDtor(VD));
640 }
641 } // namespace
642 
643 void Sema::checkAllowedCUDAInitializer(VarDecl *VD) {
644   // Return early if VD is inside a non-instantiated template function since
645   // the implicit constructor is not defined yet.
646   if (const FunctionDecl *FD =
647           dyn_cast_or_null<FunctionDecl>(VD->getDeclContext()))
648     if (FD->isDependentContext())
649       return;
650 
651   // Do not check dependent variables since the ctor/dtor/initializer are not
652   // determined. Do it after instantiation.
653   if (VD->isInvalidDecl() || !VD->hasInit() || !VD->hasGlobalStorage() ||
654       IsDependentVar(VD))
655     return;
656   const Expr *Init = VD->getInit();
657   bool IsSharedVar = VD->hasAttr<CUDASharedAttr>();
658   bool IsDeviceOrConstantVar =
659       !IsSharedVar &&
660       (VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>());
661   if (IsDeviceOrConstantVar || IsSharedVar) {
662     if (HasAllowedCUDADeviceStaticInitializer(
663             *this, VD, IsSharedVar ? CICK_Shared : CICK_DeviceOrConstant))
664       return;
665     Diag(VD->getLocation(),
666          IsSharedVar ? diag::err_shared_var_init : diag::err_dynamic_var_init)
667         << Init->getSourceRange();
668     VD->setInvalidDecl();
669   } else {
670     // This is a host-side global variable.  Check that the initializer is
671     // callable from the host side.
672     const FunctionDecl *InitFn = nullptr;
673     if (const CXXConstructExpr *CE = dyn_cast<CXXConstructExpr>(Init)) {
674       InitFn = CE->getConstructor();
675     } else if (const CallExpr *CE = dyn_cast<CallExpr>(Init)) {
676       InitFn = CE->getDirectCallee();
677     }
678     if (InitFn) {
679       CUDAFunctionTarget InitFnTarget = IdentifyCUDATarget(InitFn);
680       if (InitFnTarget != CFT_Host && InitFnTarget != CFT_HostDevice) {
681         Diag(VD->getLocation(), diag::err_ref_bad_target_global_initializer)
682             << InitFnTarget << InitFn;
683         Diag(InitFn->getLocation(), diag::note_previous_decl) << InitFn;
684         VD->setInvalidDecl();
685       }
686     }
687   }
688 }
689 
690 void Sema::CUDARecordImplicitHostDeviceFuncUsedByDevice(
691     const FunctionDecl *Callee) {
692   FunctionDecl *Caller = getCurFunctionDecl(/*AllowLambda=*/true);
693   if (!Caller)
694     return;
695 
696   if (!isCUDAImplicitHostDeviceFunction(Callee))
697     return;
698 
699   CUDAFunctionTarget CallerTarget = IdentifyCUDATarget(Caller);
700 
701   // Record whether an implicit host device function is used on device side.
702   if (CallerTarget != CFT_Device && CallerTarget != CFT_Global &&
703       (CallerTarget != CFT_HostDevice ||
704        (isCUDAImplicitHostDeviceFunction(Caller) &&
705         !getASTContext().CUDAImplicitHostDeviceFunUsedByDevice.count(Caller))))
706     return;
707 
708   getASTContext().CUDAImplicitHostDeviceFunUsedByDevice.insert(Callee);
709 }
710 
711 // With -fcuda-host-device-constexpr, an unattributed constexpr function is
712 // treated as implicitly __host__ __device__, unless:
713 //  * it is a variadic function (device-side variadic functions are not
714 //    allowed), or
715 //  * a __device__ function with this signature was already declared, in which
716 //    case in which case we output an error, unless the __device__ decl is in a
717 //    system header, in which case we leave the constexpr function unattributed.
718 //
719 // In addition, all function decls are treated as __host__ __device__ when
720 // ForceCUDAHostDeviceDepth > 0 (corresponding to code within a
721 //   #pragma clang force_cuda_host_device_begin/end
722 // pair).
723 void Sema::maybeAddCUDAHostDeviceAttrs(FunctionDecl *NewD,
724                                        const LookupResult &Previous) {
725   assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
726 
727   if (ForceCUDAHostDeviceDepth > 0) {
728     if (!NewD->hasAttr<CUDAHostAttr>())
729       NewD->addAttr(CUDAHostAttr::CreateImplicit(Context));
730     if (!NewD->hasAttr<CUDADeviceAttr>())
731       NewD->addAttr(CUDADeviceAttr::CreateImplicit(Context));
732     return;
733   }
734 
735   // If a template function has no host/device/global attributes,
736   // make it implicitly host device function.
737   if (getLangOpts().OffloadImplicitHostDeviceTemplates &&
738       !NewD->hasAttr<CUDAHostAttr>() && !NewD->hasAttr<CUDADeviceAttr>() &&
739       !NewD->hasAttr<CUDAGlobalAttr>() &&
740       (NewD->getDescribedFunctionTemplate() ||
741        NewD->isFunctionTemplateSpecialization())) {
742     NewD->addAttr(CUDAHostAttr::CreateImplicit(Context));
743     NewD->addAttr(CUDADeviceAttr::CreateImplicit(Context));
744     return;
745   }
746 
747   if (!getLangOpts().CUDAHostDeviceConstexpr || !NewD->isConstexpr() ||
748       NewD->isVariadic() || NewD->hasAttr<CUDAHostAttr>() ||
749       NewD->hasAttr<CUDADeviceAttr>() || NewD->hasAttr<CUDAGlobalAttr>())
750     return;
751 
752   // Is D a __device__ function with the same signature as NewD, ignoring CUDA
753   // attributes?
754   auto IsMatchingDeviceFn = [&](NamedDecl *D) {
755     if (UsingShadowDecl *Using = dyn_cast<UsingShadowDecl>(D))
756       D = Using->getTargetDecl();
757     FunctionDecl *OldD = D->getAsFunction();
758     return OldD && OldD->hasAttr<CUDADeviceAttr>() &&
759            !OldD->hasAttr<CUDAHostAttr>() &&
760            !IsOverload(NewD, OldD, /* UseMemberUsingDeclRules = */ false,
761                        /* ConsiderCudaAttrs = */ false);
762   };
763   auto It = llvm::find_if(Previous, IsMatchingDeviceFn);
764   if (It != Previous.end()) {
765     // We found a __device__ function with the same name and signature as NewD
766     // (ignoring CUDA attrs).  This is an error unless that function is defined
767     // in a system header, in which case we simply return without making NewD
768     // host+device.
769     NamedDecl *Match = *It;
770     if (!getSourceManager().isInSystemHeader(Match->getLocation())) {
771       Diag(NewD->getLocation(),
772            diag::err_cuda_unattributed_constexpr_cannot_overload_device)
773           << NewD;
774       Diag(Match->getLocation(),
775            diag::note_cuda_conflicting_device_function_declared_here);
776     }
777     return;
778   }
779 
780   NewD->addAttr(CUDAHostAttr::CreateImplicit(Context));
781   NewD->addAttr(CUDADeviceAttr::CreateImplicit(Context));
782 }
783 
784 // TODO: `__constant__` memory may be a limited resource for certain targets.
785 // A safeguard may be needed at the end of compilation pipeline if
786 // `__constant__` memory usage goes beyond limit.
787 void Sema::MaybeAddCUDAConstantAttr(VarDecl *VD) {
788   // Do not promote dependent variables since the cotr/dtor/initializer are
789   // not determined. Do it after instantiation.
790   if (getLangOpts().CUDAIsDevice && !VD->hasAttr<CUDAConstantAttr>() &&
791       !VD->hasAttr<CUDASharedAttr>() &&
792       (VD->isFileVarDecl() || VD->isStaticDataMember()) &&
793       !IsDependentVar(VD) &&
794       ((VD->isConstexpr() || VD->getType().isConstQualified()) &&
795        HasAllowedCUDADeviceStaticInitializer(*this, VD,
796                                              CICK_DeviceOrConstant))) {
797     VD->addAttr(CUDAConstantAttr::CreateImplicit(getASTContext()));
798   }
799 }
800 
801 Sema::SemaDiagnosticBuilder Sema::CUDADiagIfDeviceCode(SourceLocation Loc,
802                                                        unsigned DiagID) {
803   assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
804   FunctionDecl *CurFunContext = getCurFunctionDecl(/*AllowLambda=*/true);
805   SemaDiagnosticBuilder::Kind DiagKind = [&] {
806     if (!CurFunContext)
807       return SemaDiagnosticBuilder::K_Nop;
808     switch (CurrentCUDATarget()) {
809     case CFT_Global:
810     case CFT_Device:
811       return SemaDiagnosticBuilder::K_Immediate;
812     case CFT_HostDevice:
813       // An HD function counts as host code if we're compiling for host, and
814       // device code if we're compiling for device.  Defer any errors in device
815       // mode until the function is known-emitted.
816       if (!getLangOpts().CUDAIsDevice)
817         return SemaDiagnosticBuilder::K_Nop;
818       if (IsLastErrorImmediate && Diags.getDiagnosticIDs()->isBuiltinNote(DiagID))
819         return SemaDiagnosticBuilder::K_Immediate;
820       return (getEmissionStatus(CurFunContext) ==
821               FunctionEmissionStatus::Emitted)
822                  ? SemaDiagnosticBuilder::K_ImmediateWithCallStack
823                  : SemaDiagnosticBuilder::K_Deferred;
824     default:
825       return SemaDiagnosticBuilder::K_Nop;
826     }
827   }();
828   return SemaDiagnosticBuilder(DiagKind, Loc, DiagID, CurFunContext, *this);
829 }
830 
831 Sema::SemaDiagnosticBuilder Sema::CUDADiagIfHostCode(SourceLocation Loc,
832                                                      unsigned DiagID) {
833   assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
834   FunctionDecl *CurFunContext = getCurFunctionDecl(/*AllowLambda=*/true);
835   SemaDiagnosticBuilder::Kind DiagKind = [&] {
836     if (!CurFunContext)
837       return SemaDiagnosticBuilder::K_Nop;
838     switch (CurrentCUDATarget()) {
839     case CFT_Host:
840       return SemaDiagnosticBuilder::K_Immediate;
841     case CFT_HostDevice:
842       // An HD function counts as host code if we're compiling for host, and
843       // device code if we're compiling for device.  Defer any errors in device
844       // mode until the function is known-emitted.
845       if (getLangOpts().CUDAIsDevice)
846         return SemaDiagnosticBuilder::K_Nop;
847       if (IsLastErrorImmediate && Diags.getDiagnosticIDs()->isBuiltinNote(DiagID))
848         return SemaDiagnosticBuilder::K_Immediate;
849       return (getEmissionStatus(CurFunContext) ==
850               FunctionEmissionStatus::Emitted)
851                  ? SemaDiagnosticBuilder::K_ImmediateWithCallStack
852                  : SemaDiagnosticBuilder::K_Deferred;
853     default:
854       return SemaDiagnosticBuilder::K_Nop;
855     }
856   }();
857   return SemaDiagnosticBuilder(DiagKind, Loc, DiagID, CurFunContext, *this);
858 }
859 
860 bool Sema::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) {
861   assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
862   assert(Callee && "Callee may not be null.");
863 
864   const auto &ExprEvalCtx = currentEvaluationContext();
865   if (ExprEvalCtx.isUnevaluated() || ExprEvalCtx.isConstantEvaluated())
866     return true;
867 
868   // FIXME: Is bailing out early correct here?  Should we instead assume that
869   // the caller is a global initializer?
870   FunctionDecl *Caller = getCurFunctionDecl(/*AllowLambda=*/true);
871   if (!Caller)
872     return true;
873 
874   // If the caller is known-emitted, mark the callee as known-emitted.
875   // Otherwise, mark the call in our call graph so we can traverse it later.
876   bool CallerKnownEmitted =
877       getEmissionStatus(Caller) == FunctionEmissionStatus::Emitted;
878   SemaDiagnosticBuilder::Kind DiagKind = [this, Caller, Callee,
879                                           CallerKnownEmitted] {
880     switch (IdentifyCUDAPreference(Caller, Callee)) {
881     case CFP_Never:
882     case CFP_WrongSide:
883       assert(Caller && "Never/wrongSide calls require a non-null caller");
884       // If we know the caller will be emitted, we know this wrong-side call
885       // will be emitted, so it's an immediate error.  Otherwise, defer the
886       // error until we know the caller is emitted.
887       return CallerKnownEmitted
888                  ? SemaDiagnosticBuilder::K_ImmediateWithCallStack
889                  : SemaDiagnosticBuilder::K_Deferred;
890     default:
891       return SemaDiagnosticBuilder::K_Nop;
892     }
893   }();
894 
895   if (DiagKind == SemaDiagnosticBuilder::K_Nop) {
896     // For -fgpu-rdc, keep track of external kernels used by host functions.
897     if (LangOpts.CUDAIsDevice && LangOpts.GPURelocatableDeviceCode &&
898         Callee->hasAttr<CUDAGlobalAttr>() && !Callee->isDefined())
899       getASTContext().CUDAExternalDeviceDeclODRUsedByHost.insert(Callee);
900     return true;
901   }
902 
903   // Avoid emitting this error twice for the same location.  Using a hashtable
904   // like this is unfortunate, but because we must continue parsing as normal
905   // after encountering a deferred error, it's otherwise very tricky for us to
906   // ensure that we only emit this deferred error once.
907   if (!LocsWithCUDACallDiags.insert({Caller, Loc}).second)
908     return true;
909 
910   SemaDiagnosticBuilder(DiagKind, Loc, diag::err_ref_bad_target, Caller, *this)
911       << IdentifyCUDATarget(Callee) << /*function*/ 0 << Callee
912       << IdentifyCUDATarget(Caller);
913   if (!Callee->getBuiltinID())
914     SemaDiagnosticBuilder(DiagKind, Callee->getLocation(),
915                           diag::note_previous_decl, Caller, *this)
916         << Callee;
917   return DiagKind != SemaDiagnosticBuilder::K_Immediate &&
918          DiagKind != SemaDiagnosticBuilder::K_ImmediateWithCallStack;
919 }
920 
921 // Check the wrong-sided reference capture of lambda for CUDA/HIP.
922 // A lambda function may capture a stack variable by reference when it is
923 // defined and uses the capture by reference when the lambda is called. When
924 // the capture and use happen on different sides, the capture is invalid and
925 // should be diagnosed.
926 void Sema::CUDACheckLambdaCapture(CXXMethodDecl *Callee,
927                                   const sema::Capture &Capture) {
928   // In host compilation we only need to check lambda functions emitted on host
929   // side. In such lambda functions, a reference capture is invalid only
930   // if the lambda structure is populated by a device function or kernel then
931   // is passed to and called by a host function. However that is impossible,
932   // since a device function or kernel can only call a device function, also a
933   // kernel cannot pass a lambda back to a host function since we cannot
934   // define a kernel argument type which can hold the lambda before the lambda
935   // itself is defined.
936   if (!LangOpts.CUDAIsDevice)
937     return;
938 
939   // File-scope lambda can only do init captures for global variables, which
940   // results in passing by value for these global variables.
941   FunctionDecl *Caller = getCurFunctionDecl(/*AllowLambda=*/true);
942   if (!Caller)
943     return;
944 
945   // In device compilation, we only need to check lambda functions which are
946   // emitted on device side. For such lambdas, a reference capture is invalid
947   // only if the lambda structure is populated by a host function then passed
948   // to and called in a device function or kernel.
949   bool CalleeIsDevice = Callee->hasAttr<CUDADeviceAttr>();
950   bool CallerIsHost =
951       !Caller->hasAttr<CUDAGlobalAttr>() && !Caller->hasAttr<CUDADeviceAttr>();
952   bool ShouldCheck = CalleeIsDevice && CallerIsHost;
953   if (!ShouldCheck || !Capture.isReferenceCapture())
954     return;
955   auto DiagKind = SemaDiagnosticBuilder::K_Deferred;
956   if (Capture.isVariableCapture() && !getLangOpts().HIPStdPar) {
957     SemaDiagnosticBuilder(DiagKind, Capture.getLocation(),
958                           diag::err_capture_bad_target, Callee, *this)
959         << Capture.getVariable();
960   } else if (Capture.isThisCapture()) {
961     // Capture of this pointer is allowed since this pointer may be pointing to
962     // managed memory which is accessible on both device and host sides. It only
963     // results in invalid memory access if this pointer points to memory not
964     // accessible on device side.
965     SemaDiagnosticBuilder(DiagKind, Capture.getLocation(),
966                           diag::warn_maybe_capture_bad_target_this_ptr, Callee,
967                           *this);
968   }
969 }
970 
971 void Sema::CUDASetLambdaAttrs(CXXMethodDecl *Method) {
972   assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
973   if (Method->hasAttr<CUDAHostAttr>() || Method->hasAttr<CUDADeviceAttr>())
974     return;
975   Method->addAttr(CUDADeviceAttr::CreateImplicit(Context));
976   Method->addAttr(CUDAHostAttr::CreateImplicit(Context));
977 }
978 
979 void Sema::checkCUDATargetOverload(FunctionDecl *NewFD,
980                                    const LookupResult &Previous) {
981   assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
982   CUDAFunctionTarget NewTarget = IdentifyCUDATarget(NewFD);
983   for (NamedDecl *OldND : Previous) {
984     FunctionDecl *OldFD = OldND->getAsFunction();
985     if (!OldFD)
986       continue;
987 
988     CUDAFunctionTarget OldTarget = IdentifyCUDATarget(OldFD);
989     // Don't allow HD and global functions to overload other functions with the
990     // same signature.  We allow overloading based on CUDA attributes so that
991     // functions can have different implementations on the host and device, but
992     // HD/global functions "exist" in some sense on both the host and device, so
993     // should have the same implementation on both sides.
994     if (NewTarget != OldTarget &&
995         ((NewTarget == CFT_HostDevice &&
996           !(LangOpts.OffloadImplicitHostDeviceTemplates &&
997             isCUDAImplicitHostDeviceFunction(NewFD) &&
998             OldTarget == CFT_Device)) ||
999          (OldTarget == CFT_HostDevice &&
1000           !(LangOpts.OffloadImplicitHostDeviceTemplates &&
1001             isCUDAImplicitHostDeviceFunction(OldFD) &&
1002             NewTarget == CFT_Device)) ||
1003          (NewTarget == CFT_Global) || (OldTarget == CFT_Global)) &&
1004         !IsOverload(NewFD, OldFD, /* UseMemberUsingDeclRules = */ false,
1005                     /* ConsiderCudaAttrs = */ false)) {
1006       Diag(NewFD->getLocation(), diag::err_cuda_ovl_target)
1007           << NewTarget << NewFD->getDeclName() << OldTarget << OldFD;
1008       Diag(OldFD->getLocation(), diag::note_previous_declaration);
1009       NewFD->setInvalidDecl();
1010       break;
1011     }
1012   }
1013 }
1014 
1015 template <typename AttrTy>
1016 static void copyAttrIfPresent(Sema &S, FunctionDecl *FD,
1017                               const FunctionDecl &TemplateFD) {
1018   if (AttrTy *Attribute = TemplateFD.getAttr<AttrTy>()) {
1019     AttrTy *Clone = Attribute->clone(S.Context);
1020     Clone->setInherited(true);
1021     FD->addAttr(Clone);
1022   }
1023 }
1024 
1025 void Sema::inheritCUDATargetAttrs(FunctionDecl *FD,
1026                                   const FunctionTemplateDecl &TD) {
1027   const FunctionDecl &TemplateFD = *TD.getTemplatedDecl();
1028   copyAttrIfPresent<CUDAGlobalAttr>(*this, FD, TemplateFD);
1029   copyAttrIfPresent<CUDAHostAttr>(*this, FD, TemplateFD);
1030   copyAttrIfPresent<CUDADeviceAttr>(*this, FD, TemplateFD);
1031 }
1032 
1033 std::string Sema::getCudaConfigureFuncName() const {
1034   if (getLangOpts().HIP)
1035     return getLangOpts().HIPUseNewLaunchAPI ? "__hipPushCallConfiguration"
1036                                             : "hipConfigureCall";
1037 
1038   // New CUDA kernel launch sequence.
1039   if (CudaFeatureEnabled(Context.getTargetInfo().getSDKVersion(),
1040                          CudaFeature::CUDA_USES_NEW_LAUNCH))
1041     return "__cudaPushCallConfiguration";
1042 
1043   // Legacy CUDA kernel configuration call
1044   return "cudaConfigureCall";
1045 }
1046