1176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines//===--- SemaCUDA.cpp - Semantic Analysis for CUDA constructs -------------===// 2176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines// 3176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines// The LLVM Compiler Infrastructure 4176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines// 5176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines// This file is distributed under the University of Illinois Open Source 6176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines// License. See LICENSE.TXT for details. 7176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines// 8176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines//===----------------------------------------------------------------------===// 9176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines/// \file 10176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines/// \brief This file implements semantic analysis for CUDA constructs. 11176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines/// 12176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines//===----------------------------------------------------------------------===// 13176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines 14176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines#include "clang/AST/ASTContext.h" 15176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines#include "clang/AST/Decl.h" 164967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar#include "clang/AST/ExprCXX.h" 170e2c34f92f00628d48968dfea096d36381f494cbStephen Hines#include "clang/Lex/Preprocessor.h" 184967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar#include "clang/Sema/Lookup.h" 194967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar#include "clang/Sema/Sema.h" 20176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines#include "clang/Sema/SemaDiagnostic.h" 214967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar#include "clang/Sema/Template.h" 22176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines#include "llvm/ADT/Optional.h" 23176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines#include "llvm/ADT/SmallVector.h" 24176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hinesusing namespace clang; 25176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines 26176edba5311f6eff0cad2631449885ddf4fbc9eaStephen HinesExprResult Sema::ActOnCUDAExecConfigExpr(Scope *S, SourceLocation LLLLoc, 27176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines MultiExprArg ExecConfig, 28176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines SourceLocation GGGLoc) { 29176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines FunctionDecl *ConfigDecl = Context.getcudaConfigureCallDecl(); 30176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines if (!ConfigDecl) 31176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines return ExprError(Diag(LLLLoc, diag::err_undeclared_var_use) 32176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines << "cudaConfigureCall"); 33176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines QualType ConfigQTy = ConfigDecl->getType(); 34176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines 35176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines DeclRefExpr *ConfigDR = new (Context) 36176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines DeclRefExpr(ConfigDecl, false, ConfigQTy, VK_LValue, LLLLoc); 37176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines MarkFunctionReferenced(LLLLoc, ConfigDecl); 38176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines 39176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines return ActOnCallExpr(S, ConfigDR, LLLLoc, ExecConfig, GGGLoc, nullptr, 40176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines /*IsExecConfig=*/true); 41176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines} 42176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines 43176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines/// IdentifyCUDATarget - Determine the CUDA compilation target for this function 44176edba5311f6eff0cad2631449885ddf4fbc9eaStephen HinesSema::CUDAFunctionTarget Sema::IdentifyCUDATarget(const FunctionDecl *D) { 45176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines if (D->hasAttr<CUDAInvalidTargetAttr>()) 46176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines return CFT_InvalidTarget; 47176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines 48176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines if (D->hasAttr<CUDAGlobalAttr>()) 49176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines return CFT_Global; 50176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines 51176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines if (D->hasAttr<CUDADeviceAttr>()) { 52176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines if (D->hasAttr<CUDAHostAttr>()) 53176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines return CFT_HostDevice; 54176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines return CFT_Device; 55176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines } else if (D->hasAttr<CUDAHostAttr>()) { 56176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines return CFT_Host; 57176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines } else if (D->isImplicit()) { 58176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines // Some implicit declarations (like intrinsic functions) are not marked. 59176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines // Set the most lenient target on them for maximal flexibility. 60176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines return CFT_HostDevice; 61176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines } 62176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines 63176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines return CFT_Host; 64176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines} 65176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines 6687d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar// * CUDA Call preference table 6787d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar// 6887d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar// F - from, 6987d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar// T - to 7087d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar// Ph - preference in host mode 7187d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar// Pd - preference in device mode 7287d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar// H - handled in (x) 734967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar// Preferences: N:native, SS:same side, HD:host-device, WS:wrong side, --:never. 7487d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar// 754967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar// | F | T | Ph | Pd | H | 764967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar// |----+----+-----+-----+-----+ 774967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar// | d | d | N | N | (c) | 784967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar// | d | g | -- | -- | (a) | 794967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar// | d | h | -- | -- | (e) | 804967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar// | d | hd | HD | HD | (b) | 814967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar// | g | d | N | N | (c) | 824967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar// | g | g | -- | -- | (a) | 834967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar// | g | h | -- | -- | (e) | 844967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar// | g | hd | HD | HD | (b) | 854967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar// | h | d | -- | -- | (e) | 864967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar// | h | g | N | N | (c) | 874967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar// | h | h | N | N | (c) | 884967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar// | h | hd | HD | HD | (b) | 894967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar// | hd | d | WS | SS | (d) | 904967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar// | hd | g | SS | -- |(d/a)| 914967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar// | hd | h | SS | WS | (d) | 924967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar// | hd | hd | HD | HD | (b) | 9387d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar 9487d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga NainarSema::CUDAFunctionPreference 9587d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga NainarSema::IdentifyCUDAPreference(const FunctionDecl *Caller, 9687d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar const FunctionDecl *Callee) { 9787d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar assert(Callee && "Callee must be valid."); 9887d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar CUDAFunctionTarget CalleeTarget = IdentifyCUDATarget(Callee); 9987d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar CUDAFunctionTarget CallerTarget = 10087d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar (Caller != nullptr) ? IdentifyCUDATarget(Caller) : Sema::CFT_Host; 10187d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar 10287d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar // If one of the targets is invalid, the check always fails, no matter what 10387d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar // the other target is. 10487d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar if (CallerTarget == CFT_InvalidTarget || CalleeTarget == CFT_InvalidTarget) 10587d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar return CFP_Never; 10687d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar 10787d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar // (a) Can't call global from some contexts until we support CUDA's 10887d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar // dynamic parallelism. 10987d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar if (CalleeTarget == CFT_Global && 11087d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar (CallerTarget == CFT_Global || CallerTarget == CFT_Device || 11187d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar (CallerTarget == CFT_HostDevice && getLangOpts().CUDAIsDevice))) 11287d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar return CFP_Never; 11387d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar 1144967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar // (b) Calling HostDevice is OK for everyone. 1154967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar if (CalleeTarget == CFT_HostDevice) 1164967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar return CFP_HostDevice; 1174967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar 1184967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar // (c) Best case scenarios 11987d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar if (CalleeTarget == CallerTarget || 12087d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar (CallerTarget == CFT_Host && CalleeTarget == CFT_Global) || 12187d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar (CallerTarget == CFT_Global && CalleeTarget == CFT_Device)) 1224967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar return CFP_Native; 12387d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar 12487d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar // (d) HostDevice behavior depends on compilation mode. 12587d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar if (CallerTarget == CFT_HostDevice) { 1264967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar // It's OK to call a compilation-mode matching function from an HD one. 1274967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar if ((getLangOpts().CUDAIsDevice && CalleeTarget == CFT_Device) || 1284967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar (!getLangOpts().CUDAIsDevice && 1294967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar (CalleeTarget == CFT_Host || CalleeTarget == CFT_Global))) 1304967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar return CFP_SameSide; 1314967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar 1324967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar // Calls from HD to non-mode-matching functions (i.e., to host functions 1334967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar // when compiling in device mode or to device functions when compiling in 1344967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar // host mode) are allowed at the sema level, but eventually rejected if 1354967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar // they're ever codegened. TODO: Reject said calls earlier. 1364967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar return CFP_WrongSide; 13787d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar } 13887d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar 13987d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar // (e) Calling across device/host boundary is not something you should do. 14087d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar if ((CallerTarget == CFT_Host && CalleeTarget == CFT_Device) || 14187d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar (CallerTarget == CFT_Device && CalleeTarget == CFT_Host) || 14287d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar (CallerTarget == CFT_Global && CalleeTarget == CFT_Host)) 1434967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar return CFP_Never; 14487d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar 14587d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar llvm_unreachable("All cases should've been handled by now."); 14687d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar} 14787d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar 1484967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainartemplate <typename T> 1494967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainarstatic void EraseUnwantedCUDAMatchesImpl( 1504967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar Sema &S, const FunctionDecl *Caller, llvm::SmallVectorImpl<T> &Matches, 1514967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar std::function<const FunctionDecl *(const T &)> FetchDecl) { 15287d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar if (Matches.size() <= 1) 15387d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar return; 15487d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar 1554967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar // Gets the CUDA function preference for a call from Caller to Match. 1564967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar auto GetCFP = [&](const T &Match) { 1574967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar return S.IdentifyCUDAPreference(Caller, FetchDecl(Match)); 1584967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar }; 1594967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar 16087d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar // Find the best call preference among the functions in Matches. 1614967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar Sema::CUDAFunctionPreference BestCFP = GetCFP(*std::max_element( 1624967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar Matches.begin(), Matches.end(), 1634967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar [&](const T &M1, const T &M2) { return GetCFP(M1) < GetCFP(M2); })); 16487d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar 16587d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar // Erase all functions with lower priority. 1664967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar Matches.erase( 1674967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar llvm::remove_if(Matches, 1684967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar [&](const T &Match) { return GetCFP(Match) < BestCFP; }), 1694967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar Matches.end()); 17087d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar} 17187d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar 17287d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainarvoid Sema::EraseUnwantedCUDAMatches(const FunctionDecl *Caller, 17387d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar SmallVectorImpl<FunctionDecl *> &Matches){ 17487d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar EraseUnwantedCUDAMatchesImpl<FunctionDecl *>( 17587d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar *this, Caller, Matches, [](const FunctionDecl *item) { return item; }); 17687d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar} 17787d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar 17887d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainarvoid Sema::EraseUnwantedCUDAMatches(const FunctionDecl *Caller, 17987d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar SmallVectorImpl<DeclAccessPair> &Matches) { 18087d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar EraseUnwantedCUDAMatchesImpl<DeclAccessPair>( 18187d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar *this, Caller, Matches, [](const DeclAccessPair &item) { 18287d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar return dyn_cast<FunctionDecl>(item.getDecl()); 18387d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar }); 18487d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar} 18587d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar 18687d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainarvoid Sema::EraseUnwantedCUDAMatches( 18787d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar const FunctionDecl *Caller, 18887d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar SmallVectorImpl<std::pair<DeclAccessPair, FunctionDecl *>> &Matches){ 18987d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar EraseUnwantedCUDAMatchesImpl<std::pair<DeclAccessPair, FunctionDecl *>>( 19087d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar *this, Caller, Matches, 19187d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar [](const std::pair<DeclAccessPair, FunctionDecl *> &item) { 19287d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar return dyn_cast<FunctionDecl>(item.second); 19387d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar }); 19487d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar} 19587d948ecccffea9e9e37d0d053b246e2d6d6c47bPirama Arumuga Nainar 196176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines/// When an implicitly-declared special member has to invoke more than one 197176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines/// base/field special member, conflicts may occur in the targets of these 198176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines/// members. For example, if one base's member __host__ and another's is 199176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines/// __device__, it's a conflict. 200176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines/// This function figures out if the given targets \param Target1 and 201176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines/// \param Target2 conflict, and if they do not it fills in 202176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines/// \param ResolvedTarget with a target that resolves for both calls. 203176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines/// \return true if there's a conflict, false otherwise. 204176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hinesstatic bool 205176edba5311f6eff0cad2631449885ddf4fbc9eaStephen HinesresolveCalleeCUDATargetConflict(Sema::CUDAFunctionTarget Target1, 206176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines Sema::CUDAFunctionTarget Target2, 207176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines Sema::CUDAFunctionTarget *ResolvedTarget) { 2084967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar // Only free functions and static member functions may be global. 2094967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar assert(Target1 != Sema::CFT_Global); 2104967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar assert(Target2 != Sema::CFT_Global); 211176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines 212176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines if (Target1 == Sema::CFT_HostDevice) { 213176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines *ResolvedTarget = Target2; 214176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines } else if (Target2 == Sema::CFT_HostDevice) { 215176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines *ResolvedTarget = Target1; 216176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines } else if (Target1 != Target2) { 217176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines return true; 218176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines } else { 219176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines *ResolvedTarget = Target1; 220176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines } 221176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines 222176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines return false; 223176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines} 224176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines 225176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hinesbool Sema::inferCUDATargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl, 226176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines CXXSpecialMember CSM, 227176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines CXXMethodDecl *MemberDecl, 228176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines bool ConstRHS, 229176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines bool Diagnose) { 230176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines llvm::Optional<CUDAFunctionTarget> InferredTarget; 231176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines 232176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines // We're going to invoke special member lookup; mark that these special 233176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines // members are called from this one, and not from its caller. 234176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines ContextRAII MethodContext(*this, MemberDecl); 235176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines 236176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines // Look for special members in base classes that should be invoked from here. 237176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines // Infer the target of this member base on the ones it should call. 238176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines // Skip direct and indirect virtual bases for abstract classes. 239176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines llvm::SmallVector<const CXXBaseSpecifier *, 16> Bases; 240176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines for (const auto &B : ClassDecl->bases()) { 241176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines if (!B.isVirtual()) { 242176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines Bases.push_back(&B); 243176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines } 244176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines } 245176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines 246176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines if (!ClassDecl->isAbstract()) { 247176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines for (const auto &VB : ClassDecl->vbases()) { 248176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines Bases.push_back(&VB); 249176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines } 250176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines } 251176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines 252176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines for (const auto *B : Bases) { 253176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines const RecordType *BaseType = B->getType()->getAs<RecordType>(); 254176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines if (!BaseType) { 255176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines continue; 256176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines } 257176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines 258176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines CXXRecordDecl *BaseClassDecl = cast<CXXRecordDecl>(BaseType->getDecl()); 259176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines Sema::SpecialMemberOverloadResult *SMOR = 260176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines LookupSpecialMember(BaseClassDecl, CSM, 261176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines /* ConstArg */ ConstRHS, 262176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines /* VolatileArg */ false, 263176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines /* RValueThis */ false, 264176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines /* ConstThis */ false, 265176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines /* VolatileThis */ false); 266176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines 267176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines if (!SMOR || !SMOR->getMethod()) { 268176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines continue; 269176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines } 270176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines 271176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines CUDAFunctionTarget BaseMethodTarget = IdentifyCUDATarget(SMOR->getMethod()); 272176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines if (!InferredTarget.hasValue()) { 273176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines InferredTarget = BaseMethodTarget; 274176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines } else { 275176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines bool ResolutionError = resolveCalleeCUDATargetConflict( 276176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines InferredTarget.getValue(), BaseMethodTarget, 277176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines InferredTarget.getPointer()); 278176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines if (ResolutionError) { 279176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines if (Diagnose) { 280176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines Diag(ClassDecl->getLocation(), 281176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines diag::note_implicit_member_target_infer_collision) 282176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines << (unsigned)CSM << InferredTarget.getValue() << BaseMethodTarget; 283176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines } 284176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines MemberDecl->addAttr(CUDAInvalidTargetAttr::CreateImplicit(Context)); 285176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines return true; 286176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines } 287176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines } 288176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines } 289176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines 290176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines // Same as for bases, but now for special members of fields. 291176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines for (const auto *F : ClassDecl->fields()) { 292176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines if (F->isInvalidDecl()) { 293176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines continue; 294176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines } 295176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines 296176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines const RecordType *FieldType = 297176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines Context.getBaseElementType(F->getType())->getAs<RecordType>(); 298176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines if (!FieldType) { 299176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines continue; 300176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines } 301176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines 302176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines CXXRecordDecl *FieldRecDecl = cast<CXXRecordDecl>(FieldType->getDecl()); 303176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines Sema::SpecialMemberOverloadResult *SMOR = 304176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines LookupSpecialMember(FieldRecDecl, CSM, 305176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines /* ConstArg */ ConstRHS && !F->isMutable(), 306176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines /* VolatileArg */ false, 307176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines /* RValueThis */ false, 308176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines /* ConstThis */ false, 309176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines /* VolatileThis */ false); 310176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines 311176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines if (!SMOR || !SMOR->getMethod()) { 312176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines continue; 313176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines } 314176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines 315176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines CUDAFunctionTarget FieldMethodTarget = 316176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines IdentifyCUDATarget(SMOR->getMethod()); 317176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines if (!InferredTarget.hasValue()) { 318176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines InferredTarget = FieldMethodTarget; 319176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines } else { 320176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines bool ResolutionError = resolveCalleeCUDATargetConflict( 321176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines InferredTarget.getValue(), FieldMethodTarget, 322176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines InferredTarget.getPointer()); 323176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines if (ResolutionError) { 324176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines if (Diagnose) { 325176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines Diag(ClassDecl->getLocation(), 326176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines diag::note_implicit_member_target_infer_collision) 327176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines << (unsigned)CSM << InferredTarget.getValue() 328176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines << FieldMethodTarget; 329176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines } 330176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines MemberDecl->addAttr(CUDAInvalidTargetAttr::CreateImplicit(Context)); 331176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines return true; 332176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines } 333176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines } 334176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines } 335176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines 336176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines if (InferredTarget.hasValue()) { 337176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines if (InferredTarget.getValue() == CFT_Device) { 338176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines MemberDecl->addAttr(CUDADeviceAttr::CreateImplicit(Context)); 339176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines } else if (InferredTarget.getValue() == CFT_Host) { 340176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines MemberDecl->addAttr(CUDAHostAttr::CreateImplicit(Context)); 341176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines } else { 342176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines MemberDecl->addAttr(CUDADeviceAttr::CreateImplicit(Context)); 343176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines MemberDecl->addAttr(CUDAHostAttr::CreateImplicit(Context)); 344176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines } 345176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines } else { 346176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines // If no target was inferred, mark this member as __host__ __device__; 347176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines // it's the least restrictive option that can be invoked from any target. 348176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines MemberDecl->addAttr(CUDADeviceAttr::CreateImplicit(Context)); 349176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines MemberDecl->addAttr(CUDAHostAttr::CreateImplicit(Context)); 350176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines } 351176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines 352176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines return false; 353176edba5311f6eff0cad2631449885ddf4fbc9eaStephen Hines} 3544967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar 3554967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainarbool Sema::isEmptyCudaConstructor(SourceLocation Loc, CXXConstructorDecl *CD) { 3564967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar if (!CD->isDefined() && CD->isTemplateInstantiation()) 3574967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar InstantiateFunctionDefinition(Loc, CD->getFirstDecl()); 3584967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar 3594967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar // (E.2.3.1, CUDA 7.5) A constructor for a class type is considered 3604967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar // empty at a point in the translation unit, if it is either a 3614967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar // trivial constructor 3624967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar if (CD->isTrivial()) 3634967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar return true; 3644967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar 3654967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar // ... or it satisfies all of the following conditions: 3664967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar // The constructor function has been defined. 3674967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar // The constructor function has no parameters, 3684967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar // and the function body is an empty compound statement. 3694967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar if (!(CD->hasTrivialBody() && CD->getNumParams() == 0)) 3704967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar return false; 3714967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar 3724967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar // Its class has no virtual functions and no virtual base classes. 3734967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar if (CD->getParent()->isDynamicClass()) 3744967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar return false; 3754967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar 3764967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar // The only form of initializer allowed is an empty constructor. 3774967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar // This will recursively check all base classes and member initializers 3784967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar if (!llvm::all_of(CD->inits(), [&](const CXXCtorInitializer *CI) { 3794967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar if (const CXXConstructExpr *CE = 3804967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar dyn_cast<CXXConstructExpr>(CI->getInit())) 3814967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar return isEmptyCudaConstructor(Loc, CE->getConstructor()); 3824967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar return false; 3834967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar })) 3844967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar return false; 3854967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar 3864967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar return true; 3874967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar} 3884967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar 3894967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainarbool Sema::isEmptyCudaDestructor(SourceLocation Loc, CXXDestructorDecl *DD) { 3904967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar // No destructor -> no problem. 3914967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar if (!DD) 3924967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar return true; 3934967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar 3944967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar if (!DD->isDefined() && DD->isTemplateInstantiation()) 3954967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar InstantiateFunctionDefinition(Loc, DD->getFirstDecl()); 3964967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar 3974967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar // (E.2.3.1, CUDA 7.5) A destructor for a class type is considered 3984967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar // empty at a point in the translation unit, if it is either a 3994967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar // trivial constructor 4004967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar if (DD->isTrivial()) 4014967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar return true; 4024967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar 4034967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar // ... or it satisfies all of the following conditions: 4044967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar // The destructor function has been defined. 4054967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar // and the function body is an empty compound statement. 4064967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar if (!DD->hasTrivialBody()) 4074967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar return false; 4084967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar 4094967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar const CXXRecordDecl *ClassDecl = DD->getParent(); 4104967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar 4114967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar // Its class has no virtual functions and no virtual base classes. 4124967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar if (ClassDecl->isDynamicClass()) 4134967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar return false; 4144967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar 4154967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar // Only empty destructors are allowed. This will recursively check 4164967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar // destructors for all base classes... 4174967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar if (!llvm::all_of(ClassDecl->bases(), [&](const CXXBaseSpecifier &BS) { 4184967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar if (CXXRecordDecl *RD = BS.getType()->getAsCXXRecordDecl()) 4194967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar return isEmptyCudaDestructor(Loc, RD->getDestructor()); 4204967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar return true; 4214967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar })) 4224967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar return false; 4234967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar 4244967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar // ... and member fields. 4254967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar if (!llvm::all_of(ClassDecl->fields(), [&](const FieldDecl *Field) { 4264967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar if (CXXRecordDecl *RD = Field->getType() 4274967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar ->getBaseElementTypeUnsafe() 4284967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar ->getAsCXXRecordDecl()) 4294967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar return isEmptyCudaDestructor(Loc, RD->getDestructor()); 4304967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar return true; 4314967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar })) 4324967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar return false; 4334967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar 4344967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar return true; 4354967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar} 4364967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar 4374967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar// With -fcuda-host-device-constexpr, an unattributed constexpr function is 4384967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar// treated as implicitly __host__ __device__, unless: 4394967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar// * it is a variadic function (device-side variadic functions are not 4404967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar// allowed), or 4414967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar// * a __device__ function with this signature was already declared, in which 4424967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar// case in which case we output an error, unless the __device__ decl is in a 4434967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar// system header, in which case we leave the constexpr function unattributed. 4444967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainarvoid Sema::maybeAddCUDAHostDeviceAttrs(Scope *S, FunctionDecl *NewD, 4454967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar const LookupResult &Previous) { 4464967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar assert(getLangOpts().CUDA && "May be called only for CUDA compilations."); 4474967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar if (!getLangOpts().CUDAHostDeviceConstexpr || !NewD->isConstexpr() || 4484967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar NewD->isVariadic() || NewD->hasAttr<CUDAHostAttr>() || 4494967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar NewD->hasAttr<CUDADeviceAttr>() || NewD->hasAttr<CUDAGlobalAttr>()) 4504967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar return; 4514967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar 4524967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar // Is D a __device__ function with the same signature as NewD, ignoring CUDA 4534967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar // attributes? 4544967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar auto IsMatchingDeviceFn = [&](NamedDecl *D) { 4554967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar if (UsingShadowDecl *Using = dyn_cast<UsingShadowDecl>(D)) 4564967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar D = Using->getTargetDecl(); 4574967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar FunctionDecl *OldD = D->getAsFunction(); 4584967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar return OldD && OldD->hasAttr<CUDADeviceAttr>() && 4594967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar !OldD->hasAttr<CUDAHostAttr>() && 4604967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar !IsOverload(NewD, OldD, /* UseMemberUsingDeclRules = */ false, 4614967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar /* ConsiderCudaAttrs = */ false); 4624967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar }; 4634967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar auto It = llvm::find_if(Previous, IsMatchingDeviceFn); 4644967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar if (It != Previous.end()) { 4654967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar // We found a __device__ function with the same name and signature as NewD 4664967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar // (ignoring CUDA attrs). This is an error unless that function is defined 4674967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar // in a system header, in which case we simply return without making NewD 4684967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar // host+device. 4694967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar NamedDecl *Match = *It; 4704967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar if (!getSourceManager().isInSystemHeader(Match->getLocation())) { 4714967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar Diag(NewD->getLocation(), 4724967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar diag::err_cuda_unattributed_constexpr_cannot_overload_device) 4734967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar << NewD->getName(); 4744967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar Diag(Match->getLocation(), 4754967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar diag::note_cuda_conflicting_device_function_declared_here); 4764967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar } 4774967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar return; 4784967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar } 4794967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar 4804967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar NewD->addAttr(CUDAHostAttr::CreateImplicit(Context)); 4814967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar NewD->addAttr(CUDADeviceAttr::CreateImplicit(Context)); 4824967a710c84587c654b56c828382219c3937dacbPirama Arumuga Nainar} 483