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_Device) || 96 (!InDeviceMode && CalleeTarget != CFT_Host)) 97 return true; 98 } 99 100 return false; 101 } 102 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. 111 static bool 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. 119 return true; 120 } 121 122 if (Target1 == Sema::CFT_HostDevice) { 123 *ResolvedTarget = Target2; 124 } else if (Target2 == Sema::CFT_HostDevice) { 125 *ResolvedTarget = Target1; 126 } else if (Target1 != Target2) { 127 return true; 128 } else { 129 *ResolvedTarget = Target1; 130 } 131 132 return false; 133 } 134 135 bool Sema::inferCUDATargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl, 136 CXXSpecialMember CSM, 137 CXXMethodDecl *MemberDecl, 138 bool ConstRHS, 139 bool Diagnose) { 140 llvm::Optional<CUDAFunctionTarget> InferredTarget; 141 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); 145 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()) { 152 Bases.push_back(&B); 153 } 154 } 155 156 if (!ClassDecl->isAbstract()) { 157 for (const auto &VB : ClassDecl->vbases()) { 158 Bases.push_back(&VB); 159 } 160 } 161 162 for (const auto *B : Bases) { 163 const RecordType *BaseType = B->getType()->getAs<RecordType>(); 164 if (!BaseType) { 165 continue; 166 } 167 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); 176 177 if (!SMOR || !SMOR->getMethod()) { 178 continue; 179 } 180 181 CUDAFunctionTarget BaseMethodTarget = IdentifyCUDATarget(SMOR->getMethod()); 182 if (!InferredTarget.hasValue()) { 183 InferredTarget = BaseMethodTarget; 184 } else { 185 bool ResolutionError = resolveCalleeCUDATargetConflict( 186 InferredTarget.getValue(), BaseMethodTarget, 187 InferredTarget.getPointer()); 188 if (ResolutionError) { 189 if (Diagnose) { 190 Diag(ClassDecl->getLocation(), 191 diag::note_implicit_member_target_infer_collision) 192 << (unsigned)CSM << InferredTarget.getValue() << BaseMethodTarget; 193 } 194 MemberDecl->addAttr(CUDAInvalidTargetAttr::CreateImplicit(Context)); 195 return true; 196 } 197 } 198 } 199 200 // Same as for bases, but now for special members of fields. 201 for (const auto *F : ClassDecl->fields()) { 202 if (F->isInvalidDecl()) { 203 continue; 204 } 205 206 const RecordType *FieldType = 207 Context.getBaseElementType(F->getType())->getAs<RecordType>(); 208 if (!FieldType) { 209 continue; 210 } 211 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); 220 221 if (!SMOR || !SMOR->getMethod()) { 222 continue; 223 } 224 225 CUDAFunctionTarget FieldMethodTarget = 226 IdentifyCUDATarget(SMOR->getMethod()); 227 if (!InferredTarget.hasValue()) { 228 InferredTarget = FieldMethodTarget; 229 } else { 230 bool ResolutionError = resolveCalleeCUDATargetConflict( 231 InferredTarget.getValue(), FieldMethodTarget, 232 InferredTarget.getPointer()); 233 if (ResolutionError) { 234 if (Diagnose) { 235 Diag(ClassDecl->getLocation(), 236 diag::note_implicit_member_target_infer_collision) 237 << (unsigned)CSM << InferredTarget.getValue() 238 << FieldMethodTarget; 239 } 240 MemberDecl->addAttr(CUDAInvalidTargetAttr::CreateImplicit(Context)); 241 return true; 242 } 243 } 244 } 245 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)); 251 } else { 252 MemberDecl->addAttr(CUDADeviceAttr::CreateImplicit(Context)); 253 MemberDecl->addAttr(CUDAHostAttr::CreateImplicit(Context)); 254 } 255 } else { 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)); 260 } 261 262 return false; 263 } 264