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