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/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;
22 
23 ExprResult Sema::ActOnCUDAExecConfigExpr(Scope *S, SourceLocation LLLLoc,
24                                          MultiExprArg ExecConfig,
25                                          SourceLocation GGGLoc) {
26   FunctionDecl *ConfigDecl = Context.getcudaConfigureCallDecl();
27   if (!ConfigDecl)
28     return ExprError(Diag(LLLLoc, diag::err_undeclared_var_use)
29                      << "cudaConfigureCall");
30   QualType ConfigQTy = ConfigDecl->getType();
31 
32   DeclRefExpr *ConfigDR = new (Context)
33       DeclRefExpr(ConfigDecl, false, ConfigQTy, VK_LValue, LLLLoc);
34   MarkFunctionReferenced(LLLLoc, ConfigDecl);
35 
36   return ActOnCallExpr(S, ConfigDR, LLLLoc, ExecConfig, GGGLoc, nullptr,
37                        /*IsExecConfig=*/true);
38 }
39 
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;
44 
45   if (D->hasAttr<CUDAGlobalAttr>())
46     return CFT_Global;
47 
48   if (D->hasAttr<CUDADeviceAttr>()) {
49     if (D->hasAttr<CUDAHostAttr>())
50       return CFT_HostDevice;
51     return CFT_Device;
52   } else if (D->hasAttr<CUDAHostAttr>()) {
53     return CFT_Host;
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;
58   }
59 
60   return CFT_Host;
61 }
62 
63 bool Sema::CheckCUDATarget(const FunctionDecl *Caller,
64                            const FunctionDecl *Callee) {
65   CUDAFunctionTarget CallerTarget = IdentifyCUDATarget(Caller),
66                      CalleeTarget = IdentifyCUDATarget(Callee);
67 
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)
71     return true;
72 
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)
76     return true;
77 
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))
84     return true;
85 
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;
93 
94     bool InDeviceMode = getLangOpts().CUDAIsDevice;
95     if (!InDeviceMode && CalleeTarget != CFT_Host)
96         return true;
97     if (InDeviceMode && CalleeTarget != CFT_Device) {
98       // Allow host device functions to call host functions if explicitly
99       // requested.
100       if (CalleeTarget == CFT_Host &&
101           getLangOpts().CUDAAllowHostCallsFromHostDevice) {
102         Diag(Caller->getLocation(),
103              diag::warn_host_calls_from_host_device)
104             << Callee->getNameAsString() << Caller->getNameAsString();
105         return false;
106       }
107 
108       return true;
109     }
110   }
111 
112   return false;
113 }
114 
115 /// When an implicitly-declared special member has to invoke more than one
116 /// base/field special member, conflicts may occur in the targets of these
117 /// members. For example, if one base's member __host__ and another's is
118 /// __device__, it's a conflict.
119 /// This function figures out if the given targets \param Target1 and
120 /// \param Target2 conflict, and if they do not it fills in
121 /// \param ResolvedTarget with a target that resolves for both calls.
122 /// \return true if there's a conflict, false otherwise.
123 static bool
124 resolveCalleeCUDATargetConflict(Sema::CUDAFunctionTarget Target1,
125                                 Sema::CUDAFunctionTarget Target2,
126                                 Sema::CUDAFunctionTarget *ResolvedTarget) {
127   if (Target1 == Sema::CFT_Global && Target2 == Sema::CFT_Global) {
128     // TODO: this shouldn't happen, really. Methods cannot be marked __global__.
129     // Clang should detect this earlier and produce an error. Then this
130     // condition can be changed to an assertion.
131     return true;
132   }
133 
134   if (Target1 == Sema::CFT_HostDevice) {
135     *ResolvedTarget = Target2;
136   } else if (Target2 == Sema::CFT_HostDevice) {
137     *ResolvedTarget = Target1;
138   } else if (Target1 != Target2) {
139     return true;
140   } else {
141     *ResolvedTarget = Target1;
142   }
143 
144   return false;
145 }
146 
147 bool Sema::inferCUDATargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl,
148                                                    CXXSpecialMember CSM,
149                                                    CXXMethodDecl *MemberDecl,
150                                                    bool ConstRHS,
151                                                    bool Diagnose) {
152   llvm::Optional<CUDAFunctionTarget> InferredTarget;
153 
154   // We're going to invoke special member lookup; mark that these special
155   // members are called from this one, and not from its caller.
156   ContextRAII MethodContext(*this, MemberDecl);
157 
158   // Look for special members in base classes that should be invoked from here.
159   // Infer the target of this member base on the ones it should call.
160   // Skip direct and indirect virtual bases for abstract classes.
161   llvm::SmallVector<const CXXBaseSpecifier *, 16> Bases;
162   for (const auto &B : ClassDecl->bases()) {
163     if (!B.isVirtual()) {
164       Bases.push_back(&B);
165     }
166   }
167 
168   if (!ClassDecl->isAbstract()) {
169     for (const auto &VB : ClassDecl->vbases()) {
170       Bases.push_back(&VB);
171     }
172   }
173 
174   for (const auto *B : Bases) {
175     const RecordType *BaseType = B->getType()->getAs<RecordType>();
176     if (!BaseType) {
177       continue;
178     }
179 
180     CXXRecordDecl *BaseClassDecl = cast<CXXRecordDecl>(BaseType->getDecl());
181     Sema::SpecialMemberOverloadResult *SMOR =
182         LookupSpecialMember(BaseClassDecl, CSM,
183                             /* ConstArg */ ConstRHS,
184                             /* VolatileArg */ false,
185                             /* RValueThis */ false,
186                             /* ConstThis */ false,
187                             /* VolatileThis */ false);
188 
189     if (!SMOR || !SMOR->getMethod()) {
190       continue;
191     }
192 
193     CUDAFunctionTarget BaseMethodTarget = IdentifyCUDATarget(SMOR->getMethod());
194     if (!InferredTarget.hasValue()) {
195       InferredTarget = BaseMethodTarget;
196     } else {
197       bool ResolutionError = resolveCalleeCUDATargetConflict(
198           InferredTarget.getValue(), BaseMethodTarget,
199           InferredTarget.getPointer());
200       if (ResolutionError) {
201         if (Diagnose) {
202           Diag(ClassDecl->getLocation(),
203                diag::note_implicit_member_target_infer_collision)
204               << (unsigned)CSM << InferredTarget.getValue() << BaseMethodTarget;
205         }
206         MemberDecl->addAttr(CUDAInvalidTargetAttr::CreateImplicit(Context));
207         return true;
208       }
209     }
210   }
211 
212   // Same as for bases, but now for special members of fields.
213   for (const auto *F : ClassDecl->fields()) {
214     if (F->isInvalidDecl()) {
215       continue;
216     }
217 
218     const RecordType *FieldType =
219         Context.getBaseElementType(F->getType())->getAs<RecordType>();
220     if (!FieldType) {
221       continue;
222     }
223 
224     CXXRecordDecl *FieldRecDecl = cast<CXXRecordDecl>(FieldType->getDecl());
225     Sema::SpecialMemberOverloadResult *SMOR =
226         LookupSpecialMember(FieldRecDecl, CSM,
227                             /* ConstArg */ ConstRHS && !F->isMutable(),
228                             /* VolatileArg */ false,
229                             /* RValueThis */ false,
230                             /* ConstThis */ false,
231                             /* VolatileThis */ false);
232 
233     if (!SMOR || !SMOR->getMethod()) {
234       continue;
235     }
236 
237     CUDAFunctionTarget FieldMethodTarget =
238         IdentifyCUDATarget(SMOR->getMethod());
239     if (!InferredTarget.hasValue()) {
240       InferredTarget = FieldMethodTarget;
241     } else {
242       bool ResolutionError = resolveCalleeCUDATargetConflict(
243           InferredTarget.getValue(), FieldMethodTarget,
244           InferredTarget.getPointer());
245       if (ResolutionError) {
246         if (Diagnose) {
247           Diag(ClassDecl->getLocation(),
248                diag::note_implicit_member_target_infer_collision)
249               << (unsigned)CSM << InferredTarget.getValue()
250               << FieldMethodTarget;
251         }
252         MemberDecl->addAttr(CUDAInvalidTargetAttr::CreateImplicit(Context));
253         return true;
254       }
255     }
256   }
257 
258   if (InferredTarget.hasValue()) {
259     if (InferredTarget.getValue() == CFT_Device) {
260       MemberDecl->addAttr(CUDADeviceAttr::CreateImplicit(Context));
261     } else if (InferredTarget.getValue() == CFT_Host) {
262       MemberDecl->addAttr(CUDAHostAttr::CreateImplicit(Context));
263     } else {
264       MemberDecl->addAttr(CUDADeviceAttr::CreateImplicit(Context));
265       MemberDecl->addAttr(CUDAHostAttr::CreateImplicit(Context));
266     }
267   } else {
268     // If no target was inferred, mark this member as __host__ __device__;
269     // it's the least restrictive option that can be invoked from any target.
270     MemberDecl->addAttr(CUDADeviceAttr::CreateImplicit(Context));
271     MemberDecl->addAttr(CUDAHostAttr::CreateImplicit(Context));
272   }
273 
274   return false;
275 }
276