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/Sema/SemaCUDA.h"
14#include "clang/AST/ASTContext.h"
15#include "clang/AST/Decl.h"
16#include "clang/AST/EvaluatedExprVisitor.h"
17#include "clang/AST/ExprCXX.h"
18#include "clang/Basic/Cuda.h"
19#include "clang/Basic/TargetInfo.h"
20#include "clang/Lex/Preprocessor.h"
21#include "clang/Sema/Lookup.h"
22#include "clang/Sema/Overload.h"
23#include "clang/Sema/ScopeInfo.h"
24#include "clang/Sema/Sema.h"
25#include "clang/Sema/Template.h"
26#include "llvm/ADT/SmallVector.h"
27#include <optional>
28using namespace clang;
29
30SemaCUDA::SemaCUDA(Sema &S) : SemaBase(S) {}
31
32template <typename AttrT> static bool hasExplicitAttr(const VarDecl *D) {
33 if (!D)
34 return false;
35 if (auto *A = D->getAttr<AttrT>())
36 return !A->isImplicit();
37 return false;
38}
39
40void SemaCUDA::PushForceHostDevice() {
41 assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
42 ForceHostDeviceDepth++;
43}
44
45bool SemaCUDA::PopForceHostDevice() {
46 assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
47 if (ForceHostDeviceDepth == 0)
48 return false;
49 ForceHostDeviceDepth--;
50 return true;
51}
52
53ExprResult SemaCUDA::ActOnExecConfigExpr(Scope *S, SourceLocation LLLLoc,
54 MultiExprArg ExecConfig,
55 SourceLocation GGGLoc) {
56 bool IsDeviceKernelCall = false;
57 switch (CurrentTarget()) {
58 case CUDAFunctionTarget::Global:
59 case CUDAFunctionTarget::Device:
60 IsDeviceKernelCall = true;
61 break;
62 case CUDAFunctionTarget::HostDevice:
63 if (getLangOpts().CUDAIsDevice) {
64 IsDeviceKernelCall = true;
65 if (FunctionDecl *Caller =
66 SemaRef.getCurFunctionDecl(/*AllowLambda=*/true);
67 Caller && isImplicitHostDeviceFunction(D: Caller)) {
68 // Under the device compilation, config call under an HD function should
69 // be treated as a device kernel call. But, for implicit HD ones (such
70 // as lambdas), need to check whether RDC is enabled or not.
71 if (!getLangOpts().GPURelocatableDeviceCode)
72 IsDeviceKernelCall = false;
73 // HIP doesn't support device-side kernel call yet. Still treat it as
74 // the host-side kernel call.
75 if (getLangOpts().HIP)
76 IsDeviceKernelCall = false;
77 }
78 }
79 break;
80 default:
81 break;
82 }
83
84 if (IsDeviceKernelCall && getLangOpts().HIP)
85 return ExprError(
86 Diag(Loc: LLLLoc, DiagID: diag::err_cuda_device_kernel_launch_not_supported));
87
88 FunctionDecl *ConfigDecl = IsDeviceKernelCall
89 ? getASTContext().getcudaLaunchDeviceDecl()
90 : getASTContext().getcudaConfigureCallDecl();
91 if (!ConfigDecl)
92 return ExprError(Diag(Loc: LLLLoc, DiagID: diag::err_undeclared_var_use)
93 << (IsDeviceKernelCall ? getLaunchDeviceFuncName()
94 : getConfigureFuncName()));
95 // Additional check on the launch function if it's a device kernel call.
96 if (IsDeviceKernelCall) {
97 auto *GetParamBuf = getASTContext().getcudaGetParameterBufferDecl();
98 if (!GetParamBuf)
99 return ExprError(Diag(Loc: LLLLoc, DiagID: diag::err_undeclared_var_use)
100 << getGetParameterBufferFuncName());
101 }
102
103 QualType ConfigQTy = ConfigDecl->getType();
104
105 DeclRefExpr *ConfigDR = new (getASTContext()) DeclRefExpr(
106 getASTContext(), ConfigDecl, false, ConfigQTy, VK_LValue, LLLLoc);
107 SemaRef.MarkFunctionReferenced(Loc: LLLLoc, Func: ConfigDecl);
108
109 if (IsDeviceKernelCall) {
110 SmallVector<Expr *> Args;
111 // Use a null pointer as the kernel function, which may not be resolvable
112 // here. For example, resolving that kernel function may need additional
113 // kernel arguments.
114 llvm::APInt Zero(SemaRef.Context.getTypeSize(T: SemaRef.Context.IntTy), 0);
115 Args.push_back(Elt: IntegerLiteral::Create(C: SemaRef.Context, V: Zero,
116 type: SemaRef.Context.IntTy, l: LLLLoc));
117 // Use a null pointer as the placeholder of the parameter buffer, which
118 // should be replaced with the actual allocation later, in the codegen.
119 Args.push_back(Elt: IntegerLiteral::Create(C: SemaRef.Context, V: Zero,
120 type: SemaRef.Context.IntTy, l: LLLLoc));
121 // Add the original config arguments.
122 llvm::append_range(C&: Args, R&: ExecConfig);
123 // Add the default blockDim if it's missing.
124 if (Args.size() < 4) {
125 llvm::APInt One(SemaRef.Context.getTypeSize(T: SemaRef.Context.IntTy), 1);
126 Args.push_back(Elt: IntegerLiteral::Create(C: SemaRef.Context, V: One,
127 type: SemaRef.Context.IntTy, l: LLLLoc));
128 }
129 // Add the default sharedMemSize if it's missing.
130 if (Args.size() < 5)
131 Args.push_back(Elt: IntegerLiteral::Create(C: SemaRef.Context, V: Zero,
132 type: SemaRef.Context.IntTy, l: LLLLoc));
133 // Add the default stream if it's missing.
134 if (Args.size() < 6)
135 Args.push_back(Elt: new (SemaRef.Context) CXXNullPtrLiteralExpr(
136 SemaRef.Context.NullPtrTy, LLLLoc));
137 return SemaRef.BuildCallExpr(S, Fn: ConfigDR, LParenLoc: LLLLoc, ArgExprs: Args, RParenLoc: GGGLoc, ExecConfig: nullptr,
138 /*IsExecConfig=*/true);
139 }
140 return SemaRef.BuildCallExpr(S, Fn: ConfigDR, LParenLoc: LLLLoc, ArgExprs: ExecConfig, RParenLoc: GGGLoc, ExecConfig: nullptr,
141 /*IsExecConfig=*/true);
142}
143
144CUDAFunctionTarget SemaCUDA::IdentifyTarget(const ParsedAttributesView &Attrs) {
145 bool HasHostAttr = false;
146 bool HasDeviceAttr = false;
147 bool HasGlobalAttr = false;
148 bool HasInvalidTargetAttr = false;
149 for (const ParsedAttr &AL : Attrs) {
150 switch (AL.getKind()) {
151 case ParsedAttr::AT_CUDAGlobal:
152 HasGlobalAttr = true;
153 break;
154 case ParsedAttr::AT_CUDAHost:
155 HasHostAttr = true;
156 break;
157 case ParsedAttr::AT_CUDADevice:
158 HasDeviceAttr = true;
159 break;
160 case ParsedAttr::AT_CUDAInvalidTarget:
161 HasInvalidTargetAttr = true;
162 break;
163 default:
164 break;
165 }
166 }
167
168 if (HasInvalidTargetAttr)
169 return CUDAFunctionTarget::InvalidTarget;
170
171 if (HasGlobalAttr)
172 return CUDAFunctionTarget::Global;
173
174 if (HasHostAttr && HasDeviceAttr)
175 return CUDAFunctionTarget::HostDevice;
176
177 if (HasDeviceAttr)
178 return CUDAFunctionTarget::Device;
179
180 return CUDAFunctionTarget::Host;
181}
182
183template <typename A>
184static bool hasAttr(const Decl *D, bool IgnoreImplicitAttr) {
185 return D->hasAttrs() && llvm::any_of(D->getAttrs(), [&](Attr *Attribute) {
186 return isa<A>(Attribute) &&
187 !(IgnoreImplicitAttr && Attribute->isImplicit());
188 });
189}
190
191SemaCUDA::CUDATargetContextRAII::CUDATargetContextRAII(
192 SemaCUDA &S_, SemaCUDA::CUDATargetContextKind K, Decl *D)
193 : S(S_) {
194 SavedCtx = S.CurCUDATargetCtx;
195 assert(K == SemaCUDA::CTCK_InitGlobalVar);
196 auto *VD = dyn_cast_or_null<VarDecl>(Val: D);
197 if (VD && VD->hasGlobalStorage() && !VD->isStaticLocal()) {
198 auto Target = CUDAFunctionTarget::Host;
199 if ((hasAttr<CUDADeviceAttr>(D: VD, /*IgnoreImplicit=*/IgnoreImplicitAttr: true) &&
200 !hasAttr<CUDAHostAttr>(D: VD, /*IgnoreImplicit=*/IgnoreImplicitAttr: true)) ||
201 hasAttr<CUDASharedAttr>(D: VD, /*IgnoreImplicit=*/IgnoreImplicitAttr: true) ||
202 hasAttr<CUDAConstantAttr>(D: VD, /*IgnoreImplicit=*/IgnoreImplicitAttr: true))
203 Target = CUDAFunctionTarget::Device;
204 S.CurCUDATargetCtx = {.Target: Target, .Kind: K, .D: VD};
205 }
206}
207
208/// IdentifyTarget - Determine the CUDA compilation target for this function
209CUDAFunctionTarget SemaCUDA::IdentifyTarget(const FunctionDecl *D,
210 bool IgnoreImplicitHDAttr) {
211 // Code that lives outside a function gets the target from CurCUDATargetCtx.
212 if (D == nullptr)
213 return CurCUDATargetCtx.Target;
214
215 // C++ deduction guides are never codegen'ed and only participate in template
216 // argument deduction. Treat them as if they were always host+device so that
217 // CUDA/HIP target checking never rejects their use based solely on target.
218 if (isa<CXXDeductionGuideDecl>(Val: D))
219 return CUDAFunctionTarget::HostDevice;
220
221 if (D->hasAttr<CUDAInvalidTargetAttr>())
222 return CUDAFunctionTarget::InvalidTarget;
223
224 if (D->hasAttr<CUDAGlobalAttr>())
225 return CUDAFunctionTarget::Global;
226
227 if (D->isConsteval())
228 return CUDAFunctionTarget::HostDevice;
229
230 if (hasAttr<CUDADeviceAttr>(D, IgnoreImplicitAttr: IgnoreImplicitHDAttr)) {
231 if (hasAttr<CUDAHostAttr>(D, IgnoreImplicitAttr: IgnoreImplicitHDAttr))
232 return CUDAFunctionTarget::HostDevice;
233 return CUDAFunctionTarget::Device;
234 } else if (hasAttr<CUDAHostAttr>(D, IgnoreImplicitAttr: IgnoreImplicitHDAttr)) {
235 return CUDAFunctionTarget::Host;
236 } else if ((D->isImplicit() || !D->isUserProvided()) &&
237 !IgnoreImplicitHDAttr) {
238 // Some implicit declarations (like intrinsic functions) are not marked.
239 // Set the most lenient target on them for maximal flexibility.
240 return CUDAFunctionTarget::HostDevice;
241 }
242
243 return CUDAFunctionTarget::Host;
244}
245
246/// IdentifyTarget - Determine the CUDA compilation target for this variable.
247SemaCUDA::CUDAVariableTarget SemaCUDA::IdentifyTarget(const VarDecl *Var) {
248 if (Var->hasAttr<HIPManagedAttr>())
249 return CVT_Unified;
250 // Only constexpr and const variabless with implicit constant attribute
251 // are emitted on both sides. Such variables are promoted to device side
252 // only if they have static constant initializers on device side.
253 if ((Var->isConstexpr() || Var->getType().isConstQualified()) &&
254 Var->hasAttr<CUDAConstantAttr>() &&
255 !hasExplicitAttr<CUDAConstantAttr>(D: Var))
256 return CVT_Both;
257 if (Var->hasAttr<CUDADeviceAttr>() || Var->hasAttr<CUDAConstantAttr>() ||
258 Var->hasAttr<CUDASharedAttr>() ||
259 Var->getType()->isCUDADeviceBuiltinSurfaceType() ||
260 Var->getType()->isCUDADeviceBuiltinTextureType())
261 return CVT_Device;
262 // Function-scope static variable without explicit device or constant
263 // attribute are emitted
264 // - on both sides in host device functions
265 // - on device side in device or global functions
266 if (auto *FD = dyn_cast<FunctionDecl>(Val: Var->getDeclContext())) {
267 switch (IdentifyTarget(D: FD)) {
268 case CUDAFunctionTarget::HostDevice:
269 return CVT_Both;
270 case CUDAFunctionTarget::Device:
271 case CUDAFunctionTarget::Global:
272 return CVT_Device;
273 default:
274 return CVT_Host;
275 }
276 }
277 return CVT_Host;
278}
279
280// * CUDA Call preference table
281//
282// F - from,
283// T - to
284// Ph - preference in host mode
285// Pd - preference in device mode
286// H - handled in (x)
287// Preferences: N:native, SS:same side, HD:host-device, WS:wrong side, --:never.
288//
289// | F | T | Ph | Pd | H |
290// |----+----+-----+-----+-----+
291// | d | d | N | N | (c) |
292// | d | g | -- | -- | (a) |
293// | d | h | -- | -- | (e) |
294// | d | hd | HD | HD | (b) |
295// | g | d | N | N | (c) |
296// | g | g | -- | -- | (a) |
297// | g | h | -- | -- | (e) |
298// | g | hd | HD | HD | (b) |
299// | h | d | -- | -- | (e) |
300// | h | g | N | N | (c) |
301// | h | h | N | N | (c) |
302// | h | hd | HD | HD | (b) |
303// | hd | d | WS | SS | (d) |
304// | hd | g | SS | -- |(d/a)|
305// | hd | h | SS | WS | (d) |
306// | hd | hd | HD | HD | (b) |
307
308SemaCUDA::CUDAFunctionPreference
309SemaCUDA::IdentifyPreference(const FunctionDecl *Caller,
310 const FunctionDecl *Callee) {
311 assert(Callee && "Callee must be valid.");
312
313 // Treat ctor/dtor as host device function in device var initializer to allow
314 // trivial ctor/dtor without device attr to be used. Non-trivial ctor/dtor
315 // will be diagnosed by checkAllowedInitializer.
316 if (Caller == nullptr && CurCUDATargetCtx.Kind == CTCK_InitGlobalVar &&
317 CurCUDATargetCtx.Target == CUDAFunctionTarget::Device &&
318 (isa<CXXConstructorDecl>(Val: Callee) || isa<CXXDestructorDecl>(Val: Callee)))
319 return CFP_HostDevice;
320
321 CUDAFunctionTarget CallerTarget = IdentifyTarget(D: Caller);
322 CUDAFunctionTarget CalleeTarget = IdentifyTarget(D: Callee);
323
324 // If one of the targets is invalid, the check always fails, no matter what
325 // the other target is.
326 if (CallerTarget == CUDAFunctionTarget::InvalidTarget ||
327 CalleeTarget == CUDAFunctionTarget::InvalidTarget)
328 return CFP_Never;
329
330 // (a) Call global from either global or device contexts is allowed as part
331 // of CUDA's dynamic parallelism support.
332 if (CalleeTarget == CUDAFunctionTarget::Global &&
333 (CallerTarget == CUDAFunctionTarget::Global ||
334 CallerTarget == CUDAFunctionTarget::Device))
335 return CFP_Native;
336
337 // (b) Calling HostDevice is OK for everyone.
338 if (CalleeTarget == CUDAFunctionTarget::HostDevice)
339 return CFP_HostDevice;
340
341 // (c) Best case scenarios
342 if (CalleeTarget == CallerTarget ||
343 (CallerTarget == CUDAFunctionTarget::Host &&
344 CalleeTarget == CUDAFunctionTarget::Global) ||
345 (CallerTarget == CUDAFunctionTarget::Global &&
346 CalleeTarget == CUDAFunctionTarget::Device))
347 return CFP_Native;
348
349 // HipStdPar mode is special, in that assessing whether a device side call to
350 // a host target is deferred to a subsequent pass, and cannot unambiguously be
351 // adjudicated in the AST, hence we optimistically allow them to pass here.
352 if (getLangOpts().HIPStdPar &&
353 (CallerTarget == CUDAFunctionTarget::Global ||
354 CallerTarget == CUDAFunctionTarget::Device ||
355 CallerTarget == CUDAFunctionTarget::HostDevice) &&
356 CalleeTarget == CUDAFunctionTarget::Host)
357 return CFP_HostDevice;
358
359 // (d) HostDevice behavior depends on compilation mode.
360 if (CallerTarget == CUDAFunctionTarget::HostDevice) {
361 // It's OK to call a compilation-mode matching function from an HD one.
362 if ((getLangOpts().CUDAIsDevice &&
363 (CalleeTarget == CUDAFunctionTarget::Device ||
364 CalleeTarget == CUDAFunctionTarget::Global)) ||
365 (!getLangOpts().CUDAIsDevice &&
366 (CalleeTarget == CUDAFunctionTarget::Host ||
367 CalleeTarget == CUDAFunctionTarget::Global)))
368 return CFP_SameSide;
369
370 // Calls from HD to non-mode-matching functions (i.e., to host functions
371 // when compiling in device mode or to device functions when compiling in
372 // host mode) are allowed at the sema level, but eventually rejected if
373 // they're ever codegened. TODO: Reject said calls earlier.
374 return CFP_WrongSide;
375 }
376
377 // (e) Calling across device/host boundary is not something you should do.
378 if ((CallerTarget == CUDAFunctionTarget::Host &&
379 CalleeTarget == CUDAFunctionTarget::Device) ||
380 (CallerTarget == CUDAFunctionTarget::Device &&
381 CalleeTarget == CUDAFunctionTarget::Host) ||
382 (CallerTarget == CUDAFunctionTarget::Global &&
383 CalleeTarget == CUDAFunctionTarget::Host))
384 return CFP_Never;
385
386 llvm_unreachable("All cases should've been handled by now.");
387}
388
389template <typename AttrT> static bool hasImplicitAttr(const FunctionDecl *D) {
390 if (!D)
391 return false;
392 if (auto *A = D->getAttr<AttrT>())
393 return A->isImplicit();
394 return D->isImplicit();
395}
396
397bool SemaCUDA::isImplicitHostDeviceFunction(const FunctionDecl *D) {
398 bool IsImplicitDevAttr = hasImplicitAttr<CUDADeviceAttr>(D);
399 bool IsImplicitHostAttr = hasImplicitAttr<CUDAHostAttr>(D);
400 return IsImplicitDevAttr && IsImplicitHostAttr;
401}
402
403bool SemaCUDA::isImplicitHDExplicitInstantiation(const FunctionDecl *FD) {
404 return FD && FD->isImplicitHDExplicitInstantiation();
405}
406
407void SemaCUDA::EraseUnwantedMatches(
408 const FunctionDecl *Caller,
409 SmallVectorImpl<std::pair<DeclAccessPair, FunctionDecl *>> &Matches) {
410 if (Matches.size() <= 1)
411 return;
412
413 using Pair = std::pair<DeclAccessPair, FunctionDecl *>;
414
415 // Gets the CUDA function preference for a call from Caller to Match.
416 auto GetCFP = [&](const Pair &Match) {
417 return IdentifyPreference(Caller, Callee: Match.second);
418 };
419
420 // Find the best call preference among the functions in Matches.
421 CUDAFunctionPreference BestCFP =
422 GetCFP(*llvm::max_element(Range&: Matches, C: [&](const Pair &M1, const Pair &M2) {
423 return GetCFP(M1) < GetCFP(M2);
424 }));
425
426 // Erase all functions with lower priority.
427 llvm::erase_if(C&: Matches,
428 P: [&](const Pair &Match) { return GetCFP(Match) < BestCFP; });
429}
430
431/// When an implicitly-declared special member has to invoke more than one
432/// base/field special member, conflicts may occur in the targets of these
433/// members. For example, if one base's member __host__ and another's is
434/// __device__, it's a conflict.
435/// This function figures out if the given targets \param Target1 and
436/// \param Target2 conflict, and if they do not it fills in
437/// \param ResolvedTarget with a target that resolves for both calls.
438/// \return true if there's a conflict, false otherwise.
439static bool
440resolveCalleeCUDATargetConflict(CUDAFunctionTarget Target1,
441 CUDAFunctionTarget Target2,
442 CUDAFunctionTarget *ResolvedTarget) {
443 // Only free functions and static member functions may be global.
444 assert(Target1 != CUDAFunctionTarget::Global);
445 assert(Target2 != CUDAFunctionTarget::Global);
446
447 if (Target1 == CUDAFunctionTarget::HostDevice) {
448 *ResolvedTarget = Target2;
449 } else if (Target2 == CUDAFunctionTarget::HostDevice) {
450 *ResolvedTarget = Target1;
451 } else if (Target1 != Target2) {
452 return true;
453 } else {
454 *ResolvedTarget = Target1;
455 }
456
457 return false;
458}
459
460bool SemaCUDA::inferTargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl,
461 CXXSpecialMemberKind CSM,
462 CXXMethodDecl *MemberDecl,
463 bool ConstRHS,
464 bool Diagnose) {
465 // If the defaulted special member is defined lexically outside of its
466 // owning class, or the special member already has explicit device or host
467 // attributes, do not infer.
468 bool InClass = MemberDecl->getLexicalParent() == MemberDecl->getParent();
469 bool HasH = MemberDecl->hasAttr<CUDAHostAttr>();
470 bool HasD = MemberDecl->hasAttr<CUDADeviceAttr>();
471 bool HasExplicitAttr =
472 (HasD && !MemberDecl->getAttr<CUDADeviceAttr>()->isImplicit()) ||
473 (HasH && !MemberDecl->getAttr<CUDAHostAttr>()->isImplicit());
474 if (!InClass || HasExplicitAttr)
475 return false;
476
477 std::optional<CUDAFunctionTarget> InferredTarget;
478
479 // We're going to invoke special member lookup; mark that these special
480 // members are called from this one, and not from its caller.
481 Sema::ContextRAII MethodContext(SemaRef, MemberDecl);
482
483 // Look for special members in base classes that should be invoked from here.
484 // Infer the target of this member base on the ones it should call.
485 // Skip direct and indirect virtual bases for abstract classes, except for
486 // destructors — the complete destructor variant destroys virtual bases
487 // regardless of whether the class is abstract.
488 llvm::SmallVector<const CXXBaseSpecifier *, 16> Bases;
489 for (const auto &B : ClassDecl->bases()) {
490 if (!B.isVirtual()) {
491 Bases.push_back(Elt: &B);
492 }
493 }
494
495 if (!ClassDecl->isAbstract() || CSM == CXXSpecialMemberKind::Destructor)
496 llvm::append_range(C&: Bases, R: llvm::make_pointer_range(Range: ClassDecl->vbases()));
497
498 for (const auto *B : Bases) {
499 auto *BaseClassDecl = B->getType()->getAsCXXRecordDecl();
500 if (!BaseClassDecl)
501 continue;
502
503 Sema::SpecialMemberOverloadResult SMOR =
504 SemaRef.LookupSpecialMember(D: BaseClassDecl, SM: CSM,
505 /* ConstArg */ ConstRHS,
506 /* VolatileArg */ false,
507 /* RValueThis */ false,
508 /* ConstThis */ false,
509 /* VolatileThis */ false);
510
511 if (!SMOR.getMethod())
512 continue;
513
514 CUDAFunctionTarget BaseMethodTarget = IdentifyTarget(D: SMOR.getMethod());
515
516 if (!InferredTarget) {
517 InferredTarget = BaseMethodTarget;
518 } else {
519 bool ResolutionError = resolveCalleeCUDATargetConflict(
520 Target1: *InferredTarget, Target2: BaseMethodTarget, ResolvedTarget: &*InferredTarget);
521 if (ResolutionError) {
522 if (Diagnose) {
523 Diag(Loc: ClassDecl->getLocation(),
524 DiagID: diag::note_implicit_member_target_infer_collision)
525 << (unsigned)CSM << *InferredTarget << BaseMethodTarget;
526 }
527 MemberDecl->addAttr(
528 A: CUDAInvalidTargetAttr::CreateImplicit(Ctx&: getASTContext()));
529 return true;
530 }
531 }
532 }
533
534 // Same as for bases, but now for special members of fields.
535 for (const auto *F : ClassDecl->fields()) {
536 if (F->isInvalidDecl()) {
537 continue;
538 }
539
540 auto *FieldRecDecl =
541 getASTContext().getBaseElementType(QT: F->getType())->getAsCXXRecordDecl();
542 if (!FieldRecDecl)
543 continue;
544
545 Sema::SpecialMemberOverloadResult SMOR =
546 SemaRef.LookupSpecialMember(D: FieldRecDecl, SM: CSM,
547 /* ConstArg */ ConstRHS && !F->isMutable(),
548 /* VolatileArg */ false,
549 /* RValueThis */ false,
550 /* ConstThis */ false,
551 /* VolatileThis */ false);
552
553 if (!SMOR.getMethod())
554 continue;
555
556 CUDAFunctionTarget FieldMethodTarget = IdentifyTarget(D: SMOR.getMethod());
557
558 if (!InferredTarget) {
559 InferredTarget = FieldMethodTarget;
560 } else {
561 bool ResolutionError = resolveCalleeCUDATargetConflict(
562 Target1: *InferredTarget, Target2: FieldMethodTarget, ResolvedTarget: &*InferredTarget);
563 if (ResolutionError) {
564 if (Diagnose) {
565 Diag(Loc: ClassDecl->getLocation(),
566 DiagID: diag::note_implicit_member_target_infer_collision)
567 << (unsigned)CSM << *InferredTarget << FieldMethodTarget;
568 }
569 MemberDecl->addAttr(
570 A: CUDAInvalidTargetAttr::CreateImplicit(Ctx&: getASTContext()));
571 return true;
572 }
573 }
574 }
575
576 // If no target was inferred, mark this member as __host__ __device__;
577 // it's the least restrictive option that can be invoked from any target.
578 bool NeedsH = true, NeedsD = true;
579 if (InferredTarget) {
580 if (*InferredTarget == CUDAFunctionTarget::Device)
581 NeedsH = false;
582 else if (*InferredTarget == CUDAFunctionTarget::Host)
583 NeedsD = false;
584 }
585
586 // We either setting attributes first time, or the inferred ones must match
587 // previously set ones.
588 if (NeedsD && !HasD)
589 MemberDecl->addAttr(A: CUDADeviceAttr::CreateImplicit(Ctx&: getASTContext()));
590 if (NeedsH && !HasH)
591 MemberDecl->addAttr(A: CUDAHostAttr::CreateImplicit(Ctx&: getASTContext()));
592
593 return false;
594}
595
596bool SemaCUDA::isEmptyConstructor(SourceLocation Loc, CXXConstructorDecl *CD) {
597 if (!CD->isDefined() && CD->isTemplateInstantiation())
598 SemaRef.InstantiateFunctionDefinition(PointOfInstantiation: Loc, Function: CD->getFirstDecl());
599
600 // (E.2.3.1, CUDA 7.5) A constructor for a class type is considered
601 // empty at a point in the translation unit, if it is either a
602 // trivial constructor
603 if (CD->isTrivial())
604 return true;
605
606 // ... or it satisfies all of the following conditions:
607 // The constructor function has been defined.
608 // The constructor function has no parameters,
609 // and the function body is an empty compound statement.
610 if (!(CD->hasTrivialBody() && CD->getNumParams() == 0))
611 return false;
612
613 // Its class has no virtual functions and no virtual base classes.
614 if (CD->getParent()->isDynamicClass())
615 return false;
616
617 // Union ctor does not call ctors of its data members.
618 if (CD->getParent()->isUnion())
619 return true;
620
621 // The only form of initializer allowed is an empty constructor.
622 // This will recursively check all base classes and member initializers
623 if (!llvm::all_of(Range: CD->inits(), P: [&](const CXXCtorInitializer *CI) {
624 if (const CXXConstructExpr *CE =
625 dyn_cast<CXXConstructExpr>(Val: CI->getInit()))
626 return isEmptyConstructor(Loc, CD: CE->getConstructor());
627 return false;
628 }))
629 return false;
630
631 return true;
632}
633
634bool SemaCUDA::isEmptyDestructor(SourceLocation Loc, CXXDestructorDecl *DD) {
635 // No destructor -> no problem.
636 if (!DD)
637 return true;
638
639 if (!DD->isDefined() && DD->isTemplateInstantiation())
640 SemaRef.InstantiateFunctionDefinition(PointOfInstantiation: Loc, Function: DD->getFirstDecl());
641
642 // (E.2.3.1, CUDA 7.5) A destructor for a class type is considered
643 // empty at a point in the translation unit, if it is either a
644 // trivial constructor
645 if (DD->isTrivial())
646 return true;
647
648 // ... or it satisfies all of the following conditions:
649 // The destructor function has been defined.
650 // and the function body is an empty compound statement.
651 if (!DD->hasTrivialBody())
652 return false;
653
654 const CXXRecordDecl *ClassDecl = DD->getParent();
655
656 // Its class has no virtual functions and no virtual base classes.
657 if (ClassDecl->isDynamicClass())
658 return false;
659
660 // Union does not have base class and union dtor does not call dtors of its
661 // data members.
662 if (DD->getParent()->isUnion())
663 return true;
664
665 // Only empty destructors are allowed. This will recursively check
666 // destructors for all base classes...
667 if (!llvm::all_of(Range: ClassDecl->bases(), P: [&](const CXXBaseSpecifier &BS) {
668 if (CXXRecordDecl *RD = BS.getType()->getAsCXXRecordDecl())
669 return isEmptyDestructor(Loc, DD: RD->getDestructor());
670 return true;
671 }))
672 return false;
673
674 // ... and member fields.
675 if (!llvm::all_of(Range: ClassDecl->fields(), P: [&](const FieldDecl *Field) {
676 if (CXXRecordDecl *RD = Field->getType()
677 ->getBaseElementTypeUnsafe()
678 ->getAsCXXRecordDecl())
679 return isEmptyDestructor(Loc, DD: RD->getDestructor());
680 return true;
681 }))
682 return false;
683
684 return true;
685}
686
687namespace {
688enum CUDAInitializerCheckKind {
689 CICK_DeviceOrConstant, // Check initializer for device/constant variable
690 CICK_Shared, // Check initializer for shared variable
691};
692
693bool IsDependentVar(VarDecl *VD) {
694 if (VD->getType()->isDependentType())
695 return true;
696 if (const auto *Init = VD->getInit())
697 return Init->isValueDependent();
698 return false;
699}
700
701// Check whether a variable has an allowed initializer for a CUDA device side
702// variable with global storage. \p VD may be a host variable to be checked for
703// potential promotion to device side variable.
704//
705// CUDA/HIP allows only empty constructors as initializers for global
706// variables (see E.2.3.1, CUDA 7.5). The same restriction also applies to all
707// __shared__ variables whether they are local or not (they all are implicitly
708// static in CUDA). One exception is that CUDA allows constant initializers
709// for __constant__ and __device__ variables.
710bool HasAllowedCUDADeviceStaticInitializer(SemaCUDA &S, VarDecl *VD,
711 CUDAInitializerCheckKind CheckKind) {
712 assert(!VD->isInvalidDecl() && VD->hasGlobalStorage());
713 assert(!IsDependentVar(VD) && "do not check dependent var");
714 const Expr *Init = VD->getInit();
715 auto IsEmptyInit = [&](const Expr *Init) {
716 if (!Init)
717 return true;
718 if (const auto *CE = dyn_cast<CXXConstructExpr>(Val: Init)) {
719 return S.isEmptyConstructor(Loc: VD->getLocation(), CD: CE->getConstructor());
720 }
721 return false;
722 };
723 auto IsConstantInit = [&](const Expr *Init) {
724 assert(Init);
725 ASTContext::CUDAConstantEvalContextRAII EvalCtx(S.getASTContext(),
726 /*NoWronSidedVars=*/true);
727 return Init->isConstantInitializer(Ctx&: S.getASTContext(),
728 ForRef: VD->getType()->isReferenceType());
729 };
730 auto HasEmptyDtor = [&](VarDecl *VD) {
731 if (const auto *RD = VD->getType()->getAsCXXRecordDecl())
732 return S.isEmptyDestructor(Loc: VD->getLocation(), DD: RD->getDestructor());
733 return true;
734 };
735 if (CheckKind == CICK_Shared)
736 return IsEmptyInit(Init) && HasEmptyDtor(VD);
737 return S.getLangOpts().GPUAllowDeviceInit ||
738 ((IsEmptyInit(Init) || IsConstantInit(Init)) && HasEmptyDtor(VD));
739}
740} // namespace
741
742void SemaCUDA::checkAllowedInitializer(VarDecl *VD) {
743 // Return early if VD is inside a non-instantiated template function since
744 // the implicit constructor is not defined yet.
745 if (const FunctionDecl *FD =
746 dyn_cast_or_null<FunctionDecl>(Val: VD->getDeclContext());
747 FD && FD->isDependentContext())
748 return;
749
750 bool IsSharedVar = VD->hasAttr<CUDASharedAttr>();
751 bool IsDeviceOrConstantVar =
752 !IsSharedVar &&
753 (VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>());
754 if ((IsSharedVar || IsDeviceOrConstantVar) &&
755 VD->getType().getQualifiers().getAddressSpace() != LangAS::Default) {
756 Diag(Loc: VD->getLocation(), DiagID: diag::err_cuda_address_space_gpuvar);
757 VD->setInvalidDecl();
758 return;
759 }
760 // Do not check dependent variables since the ctor/dtor/initializer are not
761 // determined. Do it after instantiation.
762 if (VD->isInvalidDecl() || !VD->hasInit() || !VD->hasGlobalStorage() ||
763 IsDependentVar(VD))
764 return;
765 const Expr *Init = VD->getInit();
766 if (IsDeviceOrConstantVar || IsSharedVar) {
767 if (HasAllowedCUDADeviceStaticInitializer(
768 S&: *this, VD, CheckKind: IsSharedVar ? CICK_Shared : CICK_DeviceOrConstant))
769 return;
770 Diag(Loc: VD->getLocation(),
771 DiagID: IsSharedVar ? diag::err_shared_var_init : diag::err_dynamic_var_init)
772 << Init->getSourceRange();
773 VD->setInvalidDecl();
774 } else {
775 // This is a host-side global variable. Check that the initializer is
776 // callable from the host side.
777 const FunctionDecl *InitFn = nullptr;
778 if (const CXXConstructExpr *CE = dyn_cast<CXXConstructExpr>(Val: Init)) {
779 InitFn = CE->getConstructor();
780 } else if (const CallExpr *CE = dyn_cast<CallExpr>(Val: Init)) {
781 InitFn = CE->getDirectCallee();
782 }
783 if (InitFn) {
784 CUDAFunctionTarget InitFnTarget = IdentifyTarget(D: InitFn);
785 if (InitFnTarget != CUDAFunctionTarget::Host &&
786 InitFnTarget != CUDAFunctionTarget::HostDevice) {
787 Diag(Loc: VD->getLocation(), DiagID: diag::err_ref_bad_target_global_initializer)
788 << InitFnTarget << InitFn;
789 Diag(Loc: InitFn->getLocation(), DiagID: diag::note_previous_decl) << InitFn;
790 VD->setInvalidDecl();
791 }
792 }
793 struct GlobVarInitChecker : ConstEvaluatedExprVisitor<GlobVarInitChecker> {
794 using Base = ConstEvaluatedExprVisitor<GlobVarInitChecker>;
795 SemaCUDA &SCRef;
796 SourceLocation InitLoc;
797
798 GlobVarInitChecker(SemaCUDA &S, SourceLocation L)
799 : Base(S.getASTContext()), SCRef(S), InitLoc(L) {}
800 void VisitDeclRefExpr(const DeclRefExpr *DRE) {
801 if (auto *VarD = dyn_cast<VarDecl>(Val: DRE->getDecl());
802 VarD && VarD->hasAttr<HIPManagedAttr>()) {
803 SCRef.Diag(Loc: DRE->getLocation(),
804 DiagID: diag::err_cuda_invalid_use_of_managedvar);
805 SCRef.Diag(Loc: InitLoc, DiagID: diag::note_cuda_managed_var_in_glob_init);
806 }
807 }
808 };
809 GlobVarInitChecker Checker(*this, VD->getLocation());
810 Checker.Visit(S: Init);
811 }
812}
813
814void SemaCUDA::RecordImplicitHostDeviceFuncUsedByDevice(
815 const FunctionDecl *Callee) {
816 FunctionDecl *Caller = SemaRef.getCurFunctionDecl(/*AllowLambda=*/true);
817 if (!Caller)
818 return;
819
820 if (!isImplicitHostDeviceFunction(D: Callee))
821 return;
822
823 CUDAFunctionTarget CallerTarget = IdentifyTarget(D: Caller);
824
825 // Record whether an implicit host device function is used on device side.
826 if (CallerTarget != CUDAFunctionTarget::Device &&
827 CallerTarget != CUDAFunctionTarget::Global &&
828 (CallerTarget != CUDAFunctionTarget::HostDevice ||
829 (isImplicitHostDeviceFunction(D: Caller) &&
830 !getASTContext().CUDAImplicitHostDeviceFunUsedByDevice.count(V: Caller))))
831 return;
832
833 getASTContext().CUDAImplicitHostDeviceFunUsedByDevice.insert(V: Callee);
834}
835
836// With -fcuda-host-device-constexpr, an unattributed constexpr function is
837// treated as implicitly __host__ __device__, unless:
838// * it is a variadic function (device-side variadic functions are not
839// allowed), or
840// * a __device__ function with this signature was already declared, in which
841// case in which case we output an error, unless the __device__ decl is in a
842// system header, in which case we leave the constexpr function unattributed.
843//
844// In addition, all function decls are treated as __host__ __device__ when
845// ForceHostDeviceDepth > 0 (corresponding to code within a
846// #pragma clang force_cuda_host_device_begin/end
847// pair).
848void SemaCUDA::maybeAddHostDeviceAttrs(FunctionDecl *NewD,
849 const LookupResult &Previous) {
850 assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
851
852 if (ForceHostDeviceDepth > 0) {
853 if (!NewD->hasAttr<CUDAHostAttr>())
854 NewD->addAttr(A: CUDAHostAttr::CreateImplicit(Ctx&: getASTContext()));
855 if (!NewD->hasAttr<CUDADeviceAttr>())
856 NewD->addAttr(A: CUDADeviceAttr::CreateImplicit(Ctx&: getASTContext()));
857 return;
858 }
859
860 // If a template function has no host/device/global attributes,
861 // make it implicitly host device function.
862 if (getLangOpts().OffloadImplicitHostDeviceTemplates &&
863 !NewD->hasAttr<CUDAHostAttr>() && !NewD->hasAttr<CUDADeviceAttr>() &&
864 !NewD->hasAttr<CUDAGlobalAttr>() &&
865 (NewD->getDescribedFunctionTemplate() ||
866 NewD->isFunctionTemplateSpecialization())) {
867 NewD->addAttr(A: CUDAHostAttr::CreateImplicit(Ctx&: getASTContext()));
868 NewD->addAttr(A: CUDADeviceAttr::CreateImplicit(Ctx&: getASTContext()));
869 return;
870 }
871
872 if (!getLangOpts().CUDAHostDeviceConstexpr || !NewD->isConstexpr() ||
873 NewD->isVariadic() || NewD->hasAttr<CUDAHostAttr>() ||
874 NewD->hasAttr<CUDADeviceAttr>() || NewD->hasAttr<CUDAGlobalAttr>())
875 return;
876
877 // Is D a __device__ function with the same signature as NewD, ignoring CUDA
878 // attributes?
879 auto IsMatchingDeviceFn = [&](NamedDecl *D) {
880 if (UsingShadowDecl *Using = dyn_cast<UsingShadowDecl>(Val: D))
881 D = Using->getTargetDecl();
882 FunctionDecl *OldD = D->getAsFunction();
883 return OldD && OldD->hasAttr<CUDADeviceAttr>() &&
884 !OldD->hasAttr<CUDAHostAttr>() &&
885 !SemaRef.IsOverload(New: NewD, Old: OldD,
886 /* UseMemberUsingDeclRules = */ false,
887 /* ConsiderCudaAttrs = */ false);
888 };
889 auto It = llvm::find_if(Range: Previous, P: IsMatchingDeviceFn);
890 if (It != Previous.end()) {
891 // We found a __device__ function with the same name and signature as NewD
892 // (ignoring CUDA attrs). This is an error unless that function is defined
893 // in a system header, in which case we simply return without making NewD
894 // host+device.
895 NamedDecl *Match = *It;
896 if (!SemaRef.getSourceManager().isInSystemHeader(Loc: Match->getLocation())) {
897 Diag(Loc: NewD->getLocation(),
898 DiagID: diag::err_cuda_unattributed_constexpr_cannot_overload_device)
899 << NewD;
900 Diag(Loc: Match->getLocation(),
901 DiagID: diag::note_cuda_conflicting_device_function_declared_here);
902 }
903 return;
904 }
905
906 NewD->addAttr(A: CUDAHostAttr::CreateImplicit(Ctx&: getASTContext()));
907 NewD->addAttr(A: CUDADeviceAttr::CreateImplicit(Ctx&: getASTContext()));
908}
909
910// TODO: `__constant__` memory may be a limited resource for certain targets.
911// A safeguard may be needed at the end of compilation pipeline if
912// `__constant__` memory usage goes beyond limit.
913void SemaCUDA::MaybeAddConstantAttr(VarDecl *VD) {
914 // Do not promote dependent variables since the cotr/dtor/initializer are
915 // not determined. Do it after instantiation.
916 if (getLangOpts().CUDAIsDevice && !VD->hasAttr<CUDAConstantAttr>() &&
917 !VD->hasAttr<CUDASharedAttr>() &&
918 (VD->isFileVarDecl() || VD->isStaticDataMember()) &&
919 !IsDependentVar(VD) &&
920 ((VD->isConstexpr() || VD->getType().isConstQualified()) &&
921 HasAllowedCUDADeviceStaticInitializer(S&: *this, VD,
922 CheckKind: CICK_DeviceOrConstant))) {
923 VD->addAttr(A: CUDAConstantAttr::CreateImplicit(Ctx&: getASTContext()));
924 }
925}
926
927SemaBase::SemaDiagnosticBuilder SemaCUDA::DiagIfDeviceCode(SourceLocation Loc,
928 unsigned DiagID) {
929 assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
930 FunctionDecl *CurFunContext =
931 SemaRef.getCurFunctionDecl(/*AllowLambda=*/true);
932 SemaDiagnosticBuilder::Kind DiagKind = [&] {
933 if (!CurFunContext)
934 return SemaDiagnosticBuilder::K_Nop;
935 switch (CurrentTarget()) {
936 case CUDAFunctionTarget::Global:
937 case CUDAFunctionTarget::Device:
938 return SemaDiagnosticBuilder::K_Immediate;
939 case CUDAFunctionTarget::HostDevice:
940 // An HD function counts as host code if we're compiling for host, and
941 // device code if we're compiling for device. Defer any errors in device
942 // mode until the function is known-emitted.
943 if (!getLangOpts().CUDAIsDevice)
944 return SemaDiagnosticBuilder::K_Nop;
945 if (SemaRef.IsLastErrorImmediate &&
946 getDiagnostics().getDiagnosticIDs()->isNote(DiagID))
947 return SemaDiagnosticBuilder::K_Immediate;
948 if (isImplicitHDExplicitInstantiation(FD: CurFunContext))
949 return SemaDiagnosticBuilder::K_Deferred;
950 return (SemaRef.getEmissionStatus(Decl: CurFunContext) ==
951 Sema::FunctionEmissionStatus::Emitted)
952 ? SemaDiagnosticBuilder::K_ImmediateWithCallStack
953 : SemaDiagnosticBuilder::K_Deferred;
954 default:
955 return SemaDiagnosticBuilder::K_Nop;
956 }
957 }();
958 return SemaDiagnosticBuilder(DiagKind, Loc, DiagID, CurFunContext, SemaRef);
959}
960
961Sema::SemaDiagnosticBuilder SemaCUDA::DiagIfHostCode(SourceLocation Loc,
962 unsigned DiagID) {
963 assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
964 FunctionDecl *CurFunContext =
965 SemaRef.getCurFunctionDecl(/*AllowLambda=*/true);
966 SemaDiagnosticBuilder::Kind DiagKind = [&] {
967 if (!CurFunContext)
968 return SemaDiagnosticBuilder::K_Nop;
969 switch (CurrentTarget()) {
970 case CUDAFunctionTarget::Host:
971 return SemaDiagnosticBuilder::K_Immediate;
972 case CUDAFunctionTarget::HostDevice:
973 // An HD function counts as host code if we're compiling for host, and
974 // device code if we're compiling for device. Defer any errors in device
975 // mode until the function is known-emitted.
976 if (getLangOpts().CUDAIsDevice)
977 return SemaDiagnosticBuilder::K_Nop;
978 if (SemaRef.IsLastErrorImmediate &&
979 getDiagnostics().getDiagnosticIDs()->isNote(DiagID))
980 return SemaDiagnosticBuilder::K_Immediate;
981 return (SemaRef.getEmissionStatus(Decl: CurFunContext) ==
982 Sema::FunctionEmissionStatus::Emitted)
983 ? SemaDiagnosticBuilder::K_ImmediateWithCallStack
984 : SemaDiagnosticBuilder::K_Deferred;
985 default:
986 return SemaDiagnosticBuilder::K_Nop;
987 }
988 }();
989 return SemaDiagnosticBuilder(DiagKind, Loc, DiagID, CurFunContext, SemaRef);
990}
991
992bool SemaCUDA::CheckCall(SourceLocation Loc, FunctionDecl *Callee) {
993 assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
994 assert(Callee && "Callee may not be null.");
995
996 const auto &ExprEvalCtx = SemaRef.currentEvaluationContext();
997 if (ExprEvalCtx.isUnevaluated() || ExprEvalCtx.isConstantEvaluated() ||
998 ExprEvalCtx.isDiscardedStatementContext())
999 return true;
1000
1001 // C++ deduction guides participate in overload resolution but are not
1002 // callable functions and are never codegen'ed. Treat them as always
1003 // allowed for CUDA/HIP compatibility checking.
1004 if (isa<CXXDeductionGuideDecl>(Val: Callee))
1005 return true;
1006
1007 // FIXME: Is bailing out early correct here? Should we instead assume that
1008 // the caller is a global initializer?
1009 FunctionDecl *Caller = SemaRef.getCurFunctionDecl(/*AllowLambda=*/true);
1010 if (!Caller)
1011 return true;
1012
1013 // If the caller is known-emitted, mark the callee as known-emitted.
1014 // Otherwise, mark the call in our call graph so we can traverse it later.
1015 bool CallerKnownEmitted = SemaRef.getEmissionStatus(Decl: Caller) ==
1016 Sema::FunctionEmissionStatus::Emitted;
1017 bool CallerIsImplicitHDExplicitInst =
1018 isImplicitHDExplicitInstantiation(FD: Caller);
1019 SemaDiagnosticBuilder::Kind DiagKind = [this, Caller, Callee,
1020 CallerKnownEmitted,
1021 CallerIsImplicitHDExplicitInst] {
1022 switch (IdentifyPreference(Caller, Callee)) {
1023 case CFP_Never:
1024 case CFP_WrongSide:
1025 assert(Caller && "Never/wrongSide calls require a non-null caller");
1026 // If we know the caller will be emitted, we know this wrong-side call
1027 // will be emitted, so it's an immediate error. Otherwise, defer the
1028 // error until we know the caller is emitted.
1029 return (CallerKnownEmitted && !CallerIsImplicitHDExplicitInst)
1030 ? SemaDiagnosticBuilder::K_ImmediateWithCallStack
1031 : SemaDiagnosticBuilder::K_Deferred;
1032 default:
1033 return SemaDiagnosticBuilder::K_Nop;
1034 }
1035 }();
1036
1037 bool IsDeviceKernelCall = Callee == getASTContext().getcudaLaunchDeviceDecl();
1038 bool CallerHD = Caller && Caller->hasAttr<CUDAHostAttr>() &&
1039 Caller->hasAttr<CUDADeviceAttr>();
1040 bool CallerDiscard = SemaRef.getEmissionStatus(Decl: Caller) ==
1041 Sema::FunctionEmissionStatus::TemplateDiscarded;
1042 bool RDC = getLangOpts().GPURelocatableDeviceCode;
1043 if (IsDeviceKernelCall && !(CallerHD && CallerDiscard) && !RDC) {
1044 Diag(Loc, DiagID: diag::err_cuda_device_kernel_launch_require_rdc);
1045 return false;
1046 }
1047
1048 if (DiagKind == SemaDiagnosticBuilder::K_Nop) {
1049 // For -fgpu-rdc, keep track of external kernels used by host functions.
1050 if (getLangOpts().CUDAIsDevice && RDC &&
1051 Callee->hasAttr<CUDAGlobalAttr>() && !Callee->isDefined() &&
1052 (!Caller || (!Caller->getDescribedFunctionTemplate() &&
1053 getASTContext().GetGVALinkageForFunction(FD: Caller) ==
1054 GVA_StrongExternal)))
1055 getASTContext().CUDAExternalDeviceDeclODRUsedByHost.insert(X: Callee);
1056 return true;
1057 }
1058
1059 // Avoid emitting this error twice for the same location. Using a hashtable
1060 // like this is unfortunate, but because we must continue parsing as normal
1061 // after encountering a deferred error, it's otherwise very tricky for us to
1062 // ensure that we only emit this deferred error once.
1063 if (!LocsWithCUDACallDiags.insert(V: {.FD: Caller, .Loc: Loc}).second)
1064 return true;
1065
1066 SemaDiagnosticBuilder(DiagKind, Loc, diag::err_ref_bad_target, Caller,
1067 SemaRef)
1068 << IdentifyTarget(D: Callee) << /*function*/ 0 << Callee
1069 << IdentifyTarget(D: Caller);
1070 if (!Callee->getBuiltinID())
1071 SemaDiagnosticBuilder(DiagKind, Callee->getLocation(),
1072 diag::note_previous_decl, Caller, SemaRef)
1073 << Callee;
1074 return DiagKind != SemaDiagnosticBuilder::K_Immediate &&
1075 DiagKind != SemaDiagnosticBuilder::K_ImmediateWithCallStack;
1076}
1077
1078// Check the wrong-sided reference capture of lambda for CUDA/HIP.
1079// A lambda function may capture a stack variable by reference when it is
1080// defined and uses the capture by reference when the lambda is called. When
1081// the capture and use happen on different sides, the capture is invalid and
1082// should be diagnosed.
1083void SemaCUDA::CheckLambdaCapture(CXXMethodDecl *Callee,
1084 const sema::Capture &Capture) {
1085 // In host compilation we only need to check lambda functions emitted on host
1086 // side. In such lambda functions, a reference capture is invalid only
1087 // if the lambda structure is populated by a device function or kernel then
1088 // is passed to and called by a host function. However that is impossible,
1089 // since a device function or kernel can only call a device function, also a
1090 // kernel cannot pass a lambda back to a host function since we cannot
1091 // define a kernel argument type which can hold the lambda before the lambda
1092 // itself is defined.
1093 if (!getLangOpts().CUDAIsDevice)
1094 return;
1095
1096 // File-scope lambda can only do init captures for global variables, which
1097 // results in passing by value for these global variables.
1098 FunctionDecl *Caller = SemaRef.getCurFunctionDecl(/*AllowLambda=*/true);
1099 if (!Caller)
1100 return;
1101
1102 // In device compilation, we only need to check lambda functions which are
1103 // emitted on device side. For such lambdas, a reference capture is invalid
1104 // only if the lambda structure is populated by a host function then passed
1105 // to and called in a device function or kernel.
1106 bool CalleeIsDevice = Callee->hasAttr<CUDADeviceAttr>();
1107 bool CallerIsHost =
1108 !Caller->hasAttr<CUDAGlobalAttr>() && !Caller->hasAttr<CUDADeviceAttr>();
1109 bool ShouldCheck = CalleeIsDevice && CallerIsHost;
1110 if (!ShouldCheck || !Capture.isReferenceCapture())
1111 return;
1112 auto DiagKind = SemaDiagnosticBuilder::K_Deferred;
1113 if (Capture.isVariableCapture() && !getLangOpts().HIPStdPar) {
1114 SemaDiagnosticBuilder(DiagKind, Capture.getLocation(),
1115 diag::err_capture_bad_target, Callee, SemaRef)
1116 << Capture.getVariable();
1117 } else if (Capture.isThisCapture()) {
1118 // Capture of this pointer is allowed since this pointer may be pointing to
1119 // managed memory which is accessible on both device and host sides. It only
1120 // results in invalid memory access if this pointer points to memory not
1121 // accessible on device side.
1122 SemaDiagnosticBuilder(DiagKind, Capture.getLocation(),
1123 diag::warn_maybe_capture_bad_target_this_ptr, Callee,
1124 SemaRef);
1125 }
1126}
1127
1128void SemaCUDA::SetLambdaAttrs(CXXMethodDecl *Method) {
1129 assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
1130 if (Method->hasAttr<CUDAHostAttr>() || Method->hasAttr<CUDADeviceAttr>())
1131 return;
1132 Method->addAttr(A: CUDADeviceAttr::CreateImplicit(Ctx&: getASTContext()));
1133 Method->addAttr(A: CUDAHostAttr::CreateImplicit(Ctx&: getASTContext()));
1134}
1135
1136void SemaCUDA::checkTargetOverload(FunctionDecl *NewFD,
1137 const LookupResult &Previous) {
1138 assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
1139 CUDAFunctionTarget NewTarget = IdentifyTarget(D: NewFD);
1140 for (NamedDecl *OldND : Previous) {
1141 FunctionDecl *OldFD = OldND->getAsFunction();
1142 if (!OldFD)
1143 continue;
1144
1145 CUDAFunctionTarget OldTarget = IdentifyTarget(D: OldFD);
1146 // Don't allow HD and global functions to overload other functions with the
1147 // same signature. We allow overloading based on CUDA attributes so that
1148 // functions can have different implementations on the host and device, but
1149 // HD/global functions "exist" in some sense on both the host and device, so
1150 // should have the same implementation on both sides.
1151 if (NewTarget != OldTarget &&
1152 !SemaRef.IsOverload(New: NewFD, Old: OldFD, /* UseMemberUsingDeclRules = */ false,
1153 /* ConsiderCudaAttrs = */ false)) {
1154 if ((NewTarget == CUDAFunctionTarget::HostDevice &&
1155 !(getLangOpts().OffloadImplicitHostDeviceTemplates &&
1156 isImplicitHostDeviceFunction(D: NewFD) &&
1157 OldTarget == CUDAFunctionTarget::Device)) ||
1158 (OldTarget == CUDAFunctionTarget::HostDevice &&
1159 !(getLangOpts().OffloadImplicitHostDeviceTemplates &&
1160 isImplicitHostDeviceFunction(D: OldFD) &&
1161 NewTarget == CUDAFunctionTarget::Device)) ||
1162 (NewTarget == CUDAFunctionTarget::Global) ||
1163 (OldTarget == CUDAFunctionTarget::Global)) {
1164 Diag(Loc: NewFD->getLocation(), DiagID: diag::err_cuda_ovl_target)
1165 << NewTarget << NewFD->getDeclName() << OldTarget << OldFD;
1166 Diag(Loc: OldFD->getLocation(), DiagID: diag::note_previous_declaration);
1167 NewFD->setInvalidDecl();
1168 break;
1169 }
1170 if ((NewTarget == CUDAFunctionTarget::Host &&
1171 OldTarget == CUDAFunctionTarget::Device) ||
1172 (NewTarget == CUDAFunctionTarget::Device &&
1173 OldTarget == CUDAFunctionTarget::Host)) {
1174 Diag(Loc: NewFD->getLocation(), DiagID: diag::warn_offload_incompatible_redeclare)
1175 << NewTarget << OldTarget;
1176 Diag(Loc: OldFD->getLocation(), DiagID: diag::note_previous_declaration);
1177 }
1178 }
1179 }
1180}
1181
1182template <typename AttrTy>
1183static void copyAttrIfPresent(Sema &S, FunctionDecl *FD,
1184 const FunctionDecl &TemplateFD) {
1185 if (AttrTy *Attribute = TemplateFD.getAttr<AttrTy>()) {
1186 AttrTy *Clone = Attribute->clone(S.Context);
1187 Clone->setInherited(true);
1188 FD->addAttr(A: Clone);
1189 }
1190}
1191
1192void SemaCUDA::inheritTargetAttrs(FunctionDecl *FD,
1193 const FunctionTemplateDecl &TD) {
1194 const FunctionDecl &TemplateFD = *TD.getTemplatedDecl();
1195 copyAttrIfPresent<CUDAGlobalAttr>(S&: SemaRef, FD, TemplateFD);
1196 copyAttrIfPresent<CUDAHostAttr>(S&: SemaRef, FD, TemplateFD);
1197 copyAttrIfPresent<CUDADeviceAttr>(S&: SemaRef, FD, TemplateFD);
1198}
1199
1200std::string SemaCUDA::getConfigureFuncName() const {
1201 if (getLangOpts().OffloadViaLLVM)
1202 return "__llvmPushCallConfiguration";
1203
1204 if (getLangOpts().HIP)
1205 return getLangOpts().HIPUseNewLaunchAPI ? "__hipPushCallConfiguration"
1206 : "hipConfigureCall";
1207
1208 // New CUDA kernel launch sequence.
1209 if (CudaFeatureEnabled(getASTContext().getTargetInfo().getSDKVersion(),
1210 CudaFeature::CUDA_USES_NEW_LAUNCH))
1211 return "__cudaPushCallConfiguration";
1212
1213 // Legacy CUDA kernel configuration call
1214 return "cudaConfigureCall";
1215}
1216
1217std::string SemaCUDA::getGetParameterBufferFuncName() const {
1218 return "cudaGetParameterBuffer";
1219}
1220
1221std::string SemaCUDA::getLaunchDeviceFuncName() const {
1222 return "cudaLaunchDevice";
1223}
1224
1225// Record any local constexpr variables that are passed one way on the host
1226// and another on the device.
1227void SemaCUDA::recordPotentialODRUsedVariable(
1228 MultiExprArg Arguments, OverloadCandidateSet &Candidates) {
1229 sema::LambdaScopeInfo *LambdaInfo = SemaRef.getCurLambda();
1230 if (!LambdaInfo)
1231 return;
1232
1233 for (unsigned I = 0; I < Arguments.size(); ++I) {
1234 auto *DeclRef = dyn_cast<DeclRefExpr>(Val: Arguments[I]);
1235 if (!DeclRef)
1236 continue;
1237 auto *Variable = dyn_cast<VarDecl>(Val: DeclRef->getDecl());
1238 if (!Variable || !Variable->isLocalVarDecl() || !Variable->isConstexpr())
1239 continue;
1240
1241 bool HostByValue = false, HostByRef = false;
1242 bool DeviceByValue = false, DeviceByRef = false;
1243
1244 for (OverloadCandidate &Candidate : Candidates) {
1245 FunctionDecl *Callee = Candidate.Function;
1246 if (!Callee || I >= Callee->getNumParams())
1247 continue;
1248
1249 CUDAFunctionTarget Target = IdentifyTarget(D: Callee);
1250 if (Target == CUDAFunctionTarget::InvalidTarget ||
1251 Target == CUDAFunctionTarget::Global)
1252 continue;
1253
1254 bool CoversHost = (Target == CUDAFunctionTarget::Host ||
1255 Target == CUDAFunctionTarget::HostDevice);
1256 bool CoversDevice = (Target == CUDAFunctionTarget::Device ||
1257 Target == CUDAFunctionTarget::HostDevice);
1258
1259 bool IsRef = Callee->getParamDecl(i: I)->getType()->isReferenceType();
1260 HostByValue |= CoversHost && !IsRef;
1261 HostByRef |= CoversHost && IsRef;
1262 DeviceByValue |= CoversDevice && !IsRef;
1263 DeviceByRef |= CoversDevice && IsRef;
1264 }
1265
1266 if ((HostByValue && DeviceByRef) || (HostByRef && DeviceByValue))
1267 LambdaInfo->CUDAPotentialODRUsedVars.insert(Ptr: Variable);
1268 }
1269}
1270