1 //===--- SemaCUDA.cpp - Semantic Analysis for CUDA constructs -------------===//
3 // The LLVM Compiler Infrastructure
5 // This file is distributed under the University of Illinois Open Source
6 // License. See LICENSE.TXT for details.
8 //===----------------------------------------------------------------------===//
10 /// \brief This file implements semantic analysis for CUDA constructs.
12 //===----------------------------------------------------------------------===//
14 #include "clang/Sema/Sema.h"
15 #include "clang/AST/ASTContext.h"
16 #include "clang/AST/Decl.h"
17 #include "clang/Lex/Preprocessor.h"
18 #include "clang/Sema/SemaDiagnostic.h"
19 #include "llvm/ADT/Optional.h"
20 #include "llvm/ADT/SmallVector.h"
21 using namespace clang;
23 ExprResult Sema::ActOnCUDAExecConfigExpr(Scope *S, SourceLocation LLLLoc,
24 MultiExprArg ExecConfig,
25 SourceLocation GGGLoc) {
26 FunctionDecl *ConfigDecl = Context.getcudaConfigureCallDecl();
28 return ExprError(Diag(LLLLoc, diag::err_undeclared_var_use)
29 << "cudaConfigureCall");
30 QualType ConfigQTy = ConfigDecl->getType();
32 DeclRefExpr *ConfigDR = new (Context)
33 DeclRefExpr(ConfigDecl, false, ConfigQTy, VK_LValue, LLLLoc);
34 MarkFunctionReferenced(LLLLoc, ConfigDecl);
36 return ActOnCallExpr(S, ConfigDR, LLLLoc, ExecConfig, GGGLoc, nullptr,
37 /*IsExecConfig=*/true);
40 /// IdentifyCUDATarget - Determine the CUDA compilation target for this function
41 Sema::CUDAFunctionTarget Sema::IdentifyCUDATarget(const FunctionDecl *D) {
42 if (D->hasAttr<CUDAInvalidTargetAttr>())
43 return CFT_InvalidTarget;
45 if (D->hasAttr<CUDAGlobalAttr>())
48 if (D->hasAttr<CUDADeviceAttr>()) {
49 if (D->hasAttr<CUDAHostAttr>())
50 return CFT_HostDevice;
52 } else if (D->hasAttr<CUDAHostAttr>()) {
54 } else if (D->isImplicit()) {
55 // Some implicit declarations (like intrinsic functions) are not marked.
56 // Set the most lenient target on them for maximal flexibility.
57 return CFT_HostDevice;
63 bool Sema::CheckCUDATarget(const FunctionDecl *Caller,
64 const FunctionDecl *Callee) {
65 CUDAFunctionTarget CallerTarget = IdentifyCUDATarget(Caller),
66 CalleeTarget = IdentifyCUDATarget(Callee);
68 // If one of the targets is invalid, the check always fails, no matter what
69 // the other target is.
70 if (CallerTarget == CFT_InvalidTarget || CalleeTarget == CFT_InvalidTarget)
73 // CUDA B.1.1 "The __device__ qualifier declares a function that is [...]
74 // Callable from the device only."
75 if (CallerTarget == CFT_Host && CalleeTarget == CFT_Device)
78 // CUDA B.1.2 "The __global__ qualifier declares a function that is [...]
79 // Callable from the host only."
80 // CUDA B.1.3 "The __host__ qualifier declares a function that is [...]
81 // Callable from the host only."
82 if ((CallerTarget == CFT_Device || CallerTarget == CFT_Global) &&
83 (CalleeTarget == CFT_Host || CalleeTarget == CFT_Global))
86 // CUDA B.1.3 "The __device__ and __host__ qualifiers can be used together
87 // however, in which case the function is compiled for both the host and the
88 // device. The __CUDA_ARCH__ macro [...] can be used to differentiate code
89 // paths between host and device."
90 if (CallerTarget == CFT_HostDevice && CalleeTarget != CFT_HostDevice) {
91 // If the caller is implicit then the check always passes.
92 if (Caller->isImplicit()) return false;
94 bool InDeviceMode = getLangOpts().CUDAIsDevice;
95 if ((InDeviceMode && CalleeTarget != CFT_Device) ||
96 (!InDeviceMode && CalleeTarget != CFT_Host))
103 /// When an implicitly-declared special member has to invoke more than one
104 /// base/field special member, conflicts may occur in the targets of these
105 /// members. For example, if one base's member __host__ and another's is
106 /// __device__, it's a conflict.
107 /// This function figures out if the given targets \param Target1 and
108 /// \param Target2 conflict, and if they do not it fills in
109 /// \param ResolvedTarget with a target that resolves for both calls.
110 /// \return true if there's a conflict, false otherwise.
112 resolveCalleeCUDATargetConflict(Sema::CUDAFunctionTarget Target1,
113 Sema::CUDAFunctionTarget Target2,
114 Sema::CUDAFunctionTarget *ResolvedTarget) {
115 if (Target1 == Sema::CFT_Global && Target2 == Sema::CFT_Global) {
116 // TODO: this shouldn't happen, really. Methods cannot be marked __global__.
117 // Clang should detect this earlier and produce an error. Then this
118 // condition can be changed to an assertion.
122 if (Target1 == Sema::CFT_HostDevice) {
123 *ResolvedTarget = Target2;
124 } else if (Target2 == Sema::CFT_HostDevice) {
125 *ResolvedTarget = Target1;
126 } else if (Target1 != Target2) {
129 *ResolvedTarget = Target1;
135 bool Sema::inferCUDATargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl,
136 CXXSpecialMember CSM,
137 CXXMethodDecl *MemberDecl,
140 llvm::Optional<CUDAFunctionTarget> InferredTarget;
142 // We're going to invoke special member lookup; mark that these special
143 // members are called from this one, and not from its caller.
144 ContextRAII MethodContext(*this, MemberDecl);
146 // Look for special members in base classes that should be invoked from here.
147 // Infer the target of this member base on the ones it should call.
148 // Skip direct and indirect virtual bases for abstract classes.
149 llvm::SmallVector<const CXXBaseSpecifier *, 16> Bases;
150 for (const auto &B : ClassDecl->bases()) {
151 if (!B.isVirtual()) {
156 if (!ClassDecl->isAbstract()) {
157 for (const auto &VB : ClassDecl->vbases()) {
158 Bases.push_back(&VB);
162 for (const auto *B : Bases) {
163 const RecordType *BaseType = B->getType()->getAs<RecordType>();
168 CXXRecordDecl *BaseClassDecl = cast<CXXRecordDecl>(BaseType->getDecl());
169 Sema::SpecialMemberOverloadResult *SMOR =
170 LookupSpecialMember(BaseClassDecl, CSM,
171 /* ConstArg */ ConstRHS,
172 /* VolatileArg */ false,
173 /* RValueThis */ false,
174 /* ConstThis */ false,
175 /* VolatileThis */ false);
177 if (!SMOR || !SMOR->getMethod()) {
181 CUDAFunctionTarget BaseMethodTarget = IdentifyCUDATarget(SMOR->getMethod());
182 if (!InferredTarget.hasValue()) {
183 InferredTarget = BaseMethodTarget;
185 bool ResolutionError = resolveCalleeCUDATargetConflict(
186 InferredTarget.getValue(), BaseMethodTarget,
187 InferredTarget.getPointer());
188 if (ResolutionError) {
190 Diag(ClassDecl->getLocation(),
191 diag::note_implicit_member_target_infer_collision)
192 << (unsigned)CSM << InferredTarget.getValue() << BaseMethodTarget;
194 MemberDecl->addAttr(CUDAInvalidTargetAttr::CreateImplicit(Context));
200 // Same as for bases, but now for special members of fields.
201 for (const auto *F : ClassDecl->fields()) {
202 if (F->isInvalidDecl()) {
206 const RecordType *FieldType =
207 Context.getBaseElementType(F->getType())->getAs<RecordType>();
212 CXXRecordDecl *FieldRecDecl = cast<CXXRecordDecl>(FieldType->getDecl());
213 Sema::SpecialMemberOverloadResult *SMOR =
214 LookupSpecialMember(FieldRecDecl, CSM,
215 /* ConstArg */ ConstRHS && !F->isMutable(),
216 /* VolatileArg */ false,
217 /* RValueThis */ false,
218 /* ConstThis */ false,
219 /* VolatileThis */ false);
221 if (!SMOR || !SMOR->getMethod()) {
225 CUDAFunctionTarget FieldMethodTarget =
226 IdentifyCUDATarget(SMOR->getMethod());
227 if (!InferredTarget.hasValue()) {
228 InferredTarget = FieldMethodTarget;
230 bool ResolutionError = resolveCalleeCUDATargetConflict(
231 InferredTarget.getValue(), FieldMethodTarget,
232 InferredTarget.getPointer());
233 if (ResolutionError) {
235 Diag(ClassDecl->getLocation(),
236 diag::note_implicit_member_target_infer_collision)
237 << (unsigned)CSM << InferredTarget.getValue()
238 << FieldMethodTarget;
240 MemberDecl->addAttr(CUDAInvalidTargetAttr::CreateImplicit(Context));
246 if (InferredTarget.hasValue()) {
247 if (InferredTarget.getValue() == CFT_Device) {
248 MemberDecl->addAttr(CUDADeviceAttr::CreateImplicit(Context));
249 } else if (InferredTarget.getValue() == CFT_Host) {
250 MemberDecl->addAttr(CUDAHostAttr::CreateImplicit(Context));
252 MemberDecl->addAttr(CUDADeviceAttr::CreateImplicit(Context));
253 MemberDecl->addAttr(CUDAHostAttr::CreateImplicit(Context));
256 // If no target was inferred, mark this member as __host__ __device__;
257 // it's the least restrictive option that can be invoked from any target.
258 MemberDecl->addAttr(CUDADeviceAttr::CreateImplicit(Context));
259 MemberDecl->addAttr(CUDAHostAttr::CreateImplicit(Context));