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   }
52 
53   return CFT_Host;
54 }
55 
56 bool Sema::CheckCUDATarget(const FunctionDecl *Caller,
57                            const FunctionDecl *Callee) {
58   return CheckCUDATarget(IdentifyCUDATarget(Caller),
59                          IdentifyCUDATarget(Callee));
60 }
61 
62 bool Sema::CheckCUDATarget(CUDAFunctionTarget CallerTarget,
63                            CUDAFunctionTarget CalleeTarget) {
64   // If one of the targets is invalid, the check always fails, no matter what
65   // the other target is.
66   if (CallerTarget == CFT_InvalidTarget || CalleeTarget == CFT_InvalidTarget)
67     return true;
68 
69   // CUDA B.1.1 "The __device__ qualifier declares a function that is...
70   // Callable from the device only."
71   if (CallerTarget == CFT_Host && CalleeTarget == CFT_Device)
72     return true;
73 
74   // CUDA B.1.2 "The __global__ qualifier declares a function that is...
75   // Callable from the host only."
76   // CUDA B.1.3 "The __host__ qualifier declares a function that is...
77   // Callable from the host only."
78   if ((CallerTarget == CFT_Device || CallerTarget == CFT_Global) &&
79       (CalleeTarget == CFT_Host || CalleeTarget == CFT_Global))
80     return true;
81 
82   if (CallerTarget == CFT_HostDevice && CalleeTarget != CFT_HostDevice)
83     return true;
84 
85   return false;
86 }
87 
88 /// When an implicitly-declared special member has to invoke more than one
89 /// base/field special member, conflicts may occur in the targets of these
90 /// members. For example, if one base's member __host__ and another's is
91 /// __device__, it's a conflict.
92 /// This function figures out if the given targets \param Target1 and
93 /// \param Target2 conflict, and if they do not it fills in
94 /// \param ResolvedTarget with a target that resolves for both calls.
95 /// \return true if there's a conflict, false otherwise.
96 static bool
97 resolveCalleeCUDATargetConflict(Sema::CUDAFunctionTarget Target1,
98                                 Sema::CUDAFunctionTarget Target2,
99                                 Sema::CUDAFunctionTarget *ResolvedTarget) {
100   if (Target1 == Sema::CFT_Global && Target2 == Sema::CFT_Global) {
101     // TODO: this shouldn't happen, really. Methods cannot be marked __global__.
102     // Clang should detect this earlier and produce an error. Then this
103     // condition can be changed to an assertion.
104     return true;
105   }
106 
107   if (Target1 == Sema::CFT_HostDevice) {
108     *ResolvedTarget = Target2;
109   } else if (Target2 == Sema::CFT_HostDevice) {
110     *ResolvedTarget = Target1;
111   } else if (Target1 != Target2) {
112     return true;
113   } else {
114     *ResolvedTarget = Target1;
115   }
116 
117   return false;
118 }
119 
120 bool Sema::inferCUDATargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl,
121                                                    CXXSpecialMember CSM,
122                                                    CXXMethodDecl *MemberDecl,
123                                                    bool ConstRHS,
124                                                    bool Diagnose) {
125   llvm::Optional<CUDAFunctionTarget> InferredTarget;
126 
127   // We're going to invoke special member lookup; mark that these special
128   // members are called from this one, and not from its caller.
129   ContextRAII MethodContext(*this, MemberDecl);
130 
131   // Look for special members in base classes that should be invoked from here.
132   // Infer the target of this member base on the ones it should call.
133   // Skip direct and indirect virtual bases for abstract classes.
134   llvm::SmallVector<const CXXBaseSpecifier *, 16> Bases;
135   for (const auto &B : ClassDecl->bases()) {
136     if (!B.isVirtual()) {
137       Bases.push_back(&B);
138     }
139   }
140 
141   if (!ClassDecl->isAbstract()) {
142     for (const auto &VB : ClassDecl->vbases()) {
143       Bases.push_back(&VB);
144     }
145   }
146 
147   for (const auto *B : Bases) {
148     const RecordType *BaseType = B->getType()->getAs<RecordType>();
149     if (!BaseType) {
150       continue;
151     }
152 
153     CXXRecordDecl *BaseClassDecl = cast<CXXRecordDecl>(BaseType->getDecl());
154     Sema::SpecialMemberOverloadResult *SMOR =
155         LookupSpecialMember(BaseClassDecl, CSM,
156                             /* ConstArg */ ConstRHS,
157                             /* VolatileArg */ false,
158                             /* RValueThis */ false,
159                             /* ConstThis */ false,
160                             /* VolatileThis */ false);
161 
162     if (!SMOR || !SMOR->getMethod()) {
163       continue;
164     }
165 
166     CUDAFunctionTarget BaseMethodTarget = IdentifyCUDATarget(SMOR->getMethod());
167     if (!InferredTarget.hasValue()) {
168       InferredTarget = BaseMethodTarget;
169     } else {
170       bool ResolutionError = resolveCalleeCUDATargetConflict(
171           InferredTarget.getValue(), BaseMethodTarget,
172           InferredTarget.getPointer());
173       if (ResolutionError) {
174         if (Diagnose) {
175           Diag(ClassDecl->getLocation(),
176                diag::note_implicit_member_target_infer_collision)
177               << (unsigned)CSM << InferredTarget.getValue() << BaseMethodTarget;
178         }
179         MemberDecl->addAttr(CUDAInvalidTargetAttr::CreateImplicit(Context));
180         return true;
181       }
182     }
183   }
184 
185   // Same as for bases, but now for special members of fields.
186   for (const auto *F : ClassDecl->fields()) {
187     if (F->isInvalidDecl()) {
188       continue;
189     }
190 
191     const RecordType *FieldType =
192         Context.getBaseElementType(F->getType())->getAs<RecordType>();
193     if (!FieldType) {
194       continue;
195     }
196 
197     CXXRecordDecl *FieldRecDecl = cast<CXXRecordDecl>(FieldType->getDecl());
198     Sema::SpecialMemberOverloadResult *SMOR =
199         LookupSpecialMember(FieldRecDecl, CSM,
200                             /* ConstArg */ ConstRHS && !F->isMutable(),
201                             /* VolatileArg */ false,
202                             /* RValueThis */ false,
203                             /* ConstThis */ false,
204                             /* VolatileThis */ false);
205 
206     if (!SMOR || !SMOR->getMethod()) {
207       continue;
208     }
209 
210     CUDAFunctionTarget FieldMethodTarget =
211         IdentifyCUDATarget(SMOR->getMethod());
212     if (!InferredTarget.hasValue()) {
213       InferredTarget = FieldMethodTarget;
214     } else {
215       bool ResolutionError = resolveCalleeCUDATargetConflict(
216           InferredTarget.getValue(), FieldMethodTarget,
217           InferredTarget.getPointer());
218       if (ResolutionError) {
219         if (Diagnose) {
220           Diag(ClassDecl->getLocation(),
221                diag::note_implicit_member_target_infer_collision)
222               << (unsigned)CSM << InferredTarget.getValue()
223               << FieldMethodTarget;
224         }
225         MemberDecl->addAttr(CUDAInvalidTargetAttr::CreateImplicit(Context));
226         return true;
227       }
228     }
229   }
230 
231   if (InferredTarget.hasValue()) {
232     if (InferredTarget.getValue() == CFT_Device) {
233       MemberDecl->addAttr(CUDADeviceAttr::CreateImplicit(Context));
234     } else if (InferredTarget.getValue() == CFT_Host) {
235       MemberDecl->addAttr(CUDAHostAttr::CreateImplicit(Context));
236     } else {
237       MemberDecl->addAttr(CUDADeviceAttr::CreateImplicit(Context));
238       MemberDecl->addAttr(CUDAHostAttr::CreateImplicit(Context));
239     }
240   } else {
241     // If no target was inferred, mark this member as __host__ __device__;
242     // it's the least restrictive option that can be invoked from any target.
243     MemberDecl->addAttr(CUDADeviceAttr::CreateImplicit(Context));
244     MemberDecl->addAttr(CUDAHostAttr::CreateImplicit(Context));
245   }
246 
247   return false;
248 }
249