1 //===--- SemaCUDA.cpp - Semantic Analysis for CUDA constructs -------------===//
2 //
3 //                     The LLVM Compiler Infrastructure
4 //
5 // This file is distributed under the University of Illinois Open Source
6 // License. See LICENSE.TXT for details.
7 //
8 //===----------------------------------------------------------------------===//
9 /// \file
10 /// \brief This file implements semantic analysis for CUDA constructs.
11 ///
12 //===----------------------------------------------------------------------===//
13 
14 #include "clang/Sema/Sema.h"
15 #include "clang/AST/ASTContext.h"
16 #include "clang/AST/Decl.h"
17 #include "clang/Sema/SemaDiagnostic.h"
18 #include "llvm/ADT/Optional.h"
19 #include "llvm/ADT/SmallVector.h"
20 using namespace clang;
21 
22 ExprResult Sema::ActOnCUDAExecConfigExpr(Scope *S, SourceLocation LLLLoc,
23                                          MultiExprArg ExecConfig,
24                                          SourceLocation GGGLoc) {
25   FunctionDecl *ConfigDecl = Context.getcudaConfigureCallDecl();
26   if (!ConfigDecl)
27     return ExprError(Diag(LLLLoc, diag::err_undeclared_var_use)
28                      << "cudaConfigureCall");
29   QualType ConfigQTy = ConfigDecl->getType();
30 
31   DeclRefExpr *ConfigDR = new (Context)
32       DeclRefExpr(ConfigDecl, false, ConfigQTy, VK_LValue, LLLLoc);
33   MarkFunctionReferenced(LLLLoc, ConfigDecl);
34 
35   return ActOnCallExpr(S, ConfigDR, LLLLoc, ExecConfig, GGGLoc, nullptr,
36                        /*IsExecConfig=*/true);
37 }
38 
39 /// IdentifyCUDATarget - Determine the CUDA compilation target for this function
40 Sema::CUDAFunctionTarget Sema::IdentifyCUDATarget(const FunctionDecl *D) {
41   if (D->hasAttr<CUDAInvalidTargetAttr>())
42     return CFT_InvalidTarget;
43 
44   if (D->hasAttr<CUDAGlobalAttr>())
45     return CFT_Global;
46 
47   if (D->hasAttr<CUDADeviceAttr>()) {
48     if (D->hasAttr<CUDAHostAttr>())
49       return CFT_HostDevice;
50     return CFT_Device;
51   } else if (D->hasAttr<CUDAHostAttr>()) {
52     return CFT_Host;
53   } else if (D->isImplicit()) {
54     // Some implicit declarations (like intrinsic functions) are not marked.
55     // Set the most lenient target on them for maximal flexibility.
56     return CFT_HostDevice;
57   }
58 
59   return CFT_Host;
60 }
61 
62 bool Sema::CheckCUDATarget(const FunctionDecl *Caller,
63                            const FunctionDecl *Callee) {
64   return CheckCUDATarget(IdentifyCUDATarget(Caller),
65                          IdentifyCUDATarget(Callee));
66 }
67 
68 bool Sema::CheckCUDATarget(CUDAFunctionTarget CallerTarget,
69                            CUDAFunctionTarget CalleeTarget) {
70   // If one of the targets is invalid, the check always fails, no matter what
71   // the other target is.
72   if (CallerTarget == CFT_InvalidTarget || CalleeTarget == CFT_InvalidTarget)
73     return true;
74 
75   // CUDA B.1.1 "The __device__ qualifier declares a function that is...
76   // Callable from the device only."
77   if (CallerTarget == CFT_Host && CalleeTarget == CFT_Device)
78     return true;
79 
80   // CUDA B.1.2 "The __global__ qualifier declares a function that is...
81   // Callable from the host only."
82   // CUDA B.1.3 "The __host__ qualifier declares a function that is...
83   // Callable from the host only."
84   if ((CallerTarget == CFT_Device || CallerTarget == CFT_Global) &&
85       (CalleeTarget == CFT_Host || CalleeTarget == CFT_Global))
86     return true;
87 
88   if (CallerTarget == CFT_HostDevice && CalleeTarget != CFT_HostDevice)
89     return true;
90 
91   return false;
92 }
93 
94 /// When an implicitly-declared special member has to invoke more than one
95 /// base/field special member, conflicts may occur in the targets of these
96 /// members. For example, if one base's member __host__ and another's is
97 /// __device__, it's a conflict.
98 /// This function figures out if the given targets \param Target1 and
99 /// \param Target2 conflict, and if they do not it fills in
100 /// \param ResolvedTarget with a target that resolves for both calls.
101 /// \return true if there's a conflict, false otherwise.
102 static bool
103 resolveCalleeCUDATargetConflict(Sema::CUDAFunctionTarget Target1,
104                                 Sema::CUDAFunctionTarget Target2,
105                                 Sema::CUDAFunctionTarget *ResolvedTarget) {
106   if (Target1 == Sema::CFT_Global && Target2 == Sema::CFT_Global) {
107     // TODO: this shouldn't happen, really. Methods cannot be marked __global__.
108     // Clang should detect this earlier and produce an error. Then this
109     // condition can be changed to an assertion.
110     return true;
111   }
112 
113   if (Target1 == Sema::CFT_HostDevice) {
114     *ResolvedTarget = Target2;
115   } else if (Target2 == Sema::CFT_HostDevice) {
116     *ResolvedTarget = Target1;
117   } else if (Target1 != Target2) {
118     return true;
119   } else {
120     *ResolvedTarget = Target1;
121   }
122 
123   return false;
124 }
125 
126 bool Sema::inferCUDATargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl,
127                                                    CXXSpecialMember CSM,
128                                                    CXXMethodDecl *MemberDecl,
129                                                    bool ConstRHS,
130                                                    bool Diagnose) {
131   llvm::Optional<CUDAFunctionTarget> InferredTarget;
132 
133   // We're going to invoke special member lookup; mark that these special
134   // members are called from this one, and not from its caller.
135   ContextRAII MethodContext(*this, MemberDecl);
136 
137   // Look for special members in base classes that should be invoked from here.
138   // Infer the target of this member base on the ones it should call.
139   // Skip direct and indirect virtual bases for abstract classes.
140   llvm::SmallVector<const CXXBaseSpecifier *, 16> Bases;
141   for (const auto &B : ClassDecl->bases()) {
142     if (!B.isVirtual()) {
143       Bases.push_back(&B);
144     }
145   }
146 
147   if (!ClassDecl->isAbstract()) {
148     for (const auto &VB : ClassDecl->vbases()) {
149       Bases.push_back(&VB);
150     }
151   }
152 
153   for (const auto *B : Bases) {
154     const RecordType *BaseType = B->getType()->getAs<RecordType>();
155     if (!BaseType) {
156       continue;
157     }
158 
159     CXXRecordDecl *BaseClassDecl = cast<CXXRecordDecl>(BaseType->getDecl());
160     Sema::SpecialMemberOverloadResult *SMOR =
161         LookupSpecialMember(BaseClassDecl, CSM,
162                             /* ConstArg */ ConstRHS,
163                             /* VolatileArg */ false,
164                             /* RValueThis */ false,
165                             /* ConstThis */ false,
166                             /* VolatileThis */ false);
167 
168     if (!SMOR || !SMOR->getMethod()) {
169       continue;
170     }
171 
172     CUDAFunctionTarget BaseMethodTarget = IdentifyCUDATarget(SMOR->getMethod());
173     if (!InferredTarget.hasValue()) {
174       InferredTarget = BaseMethodTarget;
175     } else {
176       bool ResolutionError = resolveCalleeCUDATargetConflict(
177           InferredTarget.getValue(), BaseMethodTarget,
178           InferredTarget.getPointer());
179       if (ResolutionError) {
180         if (Diagnose) {
181           Diag(ClassDecl->getLocation(),
182                diag::note_implicit_member_target_infer_collision)
183               << (unsigned)CSM << InferredTarget.getValue() << BaseMethodTarget;
184         }
185         MemberDecl->addAttr(CUDAInvalidTargetAttr::CreateImplicit(Context));
186         return true;
187       }
188     }
189   }
190 
191   // Same as for bases, but now for special members of fields.
192   for (const auto *F : ClassDecl->fields()) {
193     if (F->isInvalidDecl()) {
194       continue;
195     }
196 
197     const RecordType *FieldType =
198         Context.getBaseElementType(F->getType())->getAs<RecordType>();
199     if (!FieldType) {
200       continue;
201     }
202 
203     CXXRecordDecl *FieldRecDecl = cast<CXXRecordDecl>(FieldType->getDecl());
204     Sema::SpecialMemberOverloadResult *SMOR =
205         LookupSpecialMember(FieldRecDecl, CSM,
206                             /* ConstArg */ ConstRHS && !F->isMutable(),
207                             /* VolatileArg */ false,
208                             /* RValueThis */ false,
209                             /* ConstThis */ false,
210                             /* VolatileThis */ false);
211 
212     if (!SMOR || !SMOR->getMethod()) {
213       continue;
214     }
215 
216     CUDAFunctionTarget FieldMethodTarget =
217         IdentifyCUDATarget(SMOR->getMethod());
218     if (!InferredTarget.hasValue()) {
219       InferredTarget = FieldMethodTarget;
220     } else {
221       bool ResolutionError = resolveCalleeCUDATargetConflict(
222           InferredTarget.getValue(), FieldMethodTarget,
223           InferredTarget.getPointer());
224       if (ResolutionError) {
225         if (Diagnose) {
226           Diag(ClassDecl->getLocation(),
227                diag::note_implicit_member_target_infer_collision)
228               << (unsigned)CSM << InferredTarget.getValue()
229               << FieldMethodTarget;
230         }
231         MemberDecl->addAttr(CUDAInvalidTargetAttr::CreateImplicit(Context));
232         return true;
233       }
234     }
235   }
236 
237   if (InferredTarget.hasValue()) {
238     if (InferredTarget.getValue() == CFT_Device) {
239       MemberDecl->addAttr(CUDADeviceAttr::CreateImplicit(Context));
240     } else if (InferredTarget.getValue() == CFT_Host) {
241       MemberDecl->addAttr(CUDAHostAttr::CreateImplicit(Context));
242     } else {
243       MemberDecl->addAttr(CUDADeviceAttr::CreateImplicit(Context));
244       MemberDecl->addAttr(CUDAHostAttr::CreateImplicit(Context));
245     }
246   } else {
247     // If no target was inferred, mark this member as __host__ __device__;
248     // it's the least restrictive option that can be invoked from any target.
249     MemberDecl->addAttr(CUDADeviceAttr::CreateImplicit(Context));
250     MemberDecl->addAttr(CUDAHostAttr::CreateImplicit(Context));
251   }
252 
253   return false;
254 }
255