Vendor import of clang trunk r338150:vendor/clang/clang-trunk-r338150

https://llvm.org/svn/llvm-project/cfe/trunk@338150

1 files changed, 73 insertions, 15 deletions

diff --git a/lib/Sema/SemaCUDA.cpp b/lib/Sema/SemaCUDA.cpp
index cac5f682275e..13dd8d936fd2 100644
--- a/lib/Sema/SemaCUDA.cpp
+++ b/lib/Sema/SemaCUDA.cpp
@@ -7,7 +7,7 @@
 //
 //===----------------------------------------------------------------------===//
 /// \file
-/// \brief This file implements semantic analysis for CUDA constructs.
+/// This file implements semantic analysis for CUDA constructs.
 ///
 //===----------------------------------------------------------------------===//
 
@@ -42,8 +42,9 @@ ExprResult Sema::ActOnCUDAExecConfigExpr(Scope *S, SourceLocation LLLLoc,
                                          SourceLocation GGGLoc) {
   FunctionDecl *ConfigDecl = Context.getcudaConfigureCallDecl();
   if (!ConfigDecl)
-    return ExprError(Diag(LLLLoc, diag::err_undeclared_var_use)
-                     << "cudaConfigureCall");
+    return ExprError(
+        Diag(LLLLoc, diag::err_undeclared_var_use)
+        << (getLangOpts().HIP ? "hipConfigureCall" : "cudaConfigureCall"));
   QualType ConfigQTy = ConfigDecl->getType();
 
   DeclRefExpr *ConfigDR = new (Context)
@@ -54,30 +55,31 @@ ExprResult Sema::ActOnCUDAExecConfigExpr(Scope *S, SourceLocation LLLLoc,
                        /*IsExecConfig=*/true);
 }
 
-Sema::CUDAFunctionTarget Sema::IdentifyCUDATarget(const AttributeList *Attr) {
+Sema::CUDAFunctionTarget
+Sema::IdentifyCUDATarget(const ParsedAttributesView &Attrs) {
   bool HasHostAttr = false;
   bool HasDeviceAttr = false;
   bool HasGlobalAttr = false;
   bool HasInvalidTargetAttr = false;
-  while (Attr) {
-    switch(Attr->getKind()){
-    case AttributeList::AT_CUDAGlobal:
+  for (const ParsedAttr &AL : Attrs) {
+    switch (AL.getKind()) {
+    case ParsedAttr::AT_CUDAGlobal:
       HasGlobalAttr = true;
       break;
-    case AttributeList::AT_CUDAHost:
+    case ParsedAttr::AT_CUDAHost:
       HasHostAttr = true;
       break;
-    case AttributeList::AT_CUDADevice:
+    case ParsedAttr::AT_CUDADevice:
       HasDeviceAttr = true;
       break;
-    case AttributeList::AT_CUDAInvalidTarget:
+    case ParsedAttr::AT_CUDAInvalidTarget:
       HasInvalidTargetAttr = true;
       break;
     default:
       break;
     }
-    Attr = Attr->getNext();
   }
+
   if (HasInvalidTargetAttr)
     return CFT_InvalidTarget;
 
@@ -471,6 +473,59 @@ bool Sema::isEmptyCudaDestructor(SourceLocation Loc, CXXDestructorDecl *DD) {
   return true;
 }
 
+void Sema::checkAllowedCUDAInitializer(VarDecl *VD) {
+  if (VD->isInvalidDecl() || !VD->hasInit() || !VD->hasGlobalStorage())
+    return;
+  const Expr *Init = VD->getInit();
+  if (VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>() ||
+      VD->hasAttr<CUDASharedAttr>()) {
+    assert(!VD->isStaticLocal() || VD->hasAttr<CUDASharedAttr>());
+    bool AllowedInit = false;
+    if (const CXXConstructExpr *CE = dyn_cast<CXXConstructExpr>(Init))
+      AllowedInit =
+          isEmptyCudaConstructor(VD->getLocation(), CE->getConstructor());
+    // We'll allow constant initializers even if it's a non-empty
+    // constructor according to CUDA rules. This deviates from NVCC,
+    // but allows us to handle things like constexpr constructors.
+    if (!AllowedInit &&
+        (VD->hasAttr<CUDADeviceAttr>() || VD->hasAttr<CUDAConstantAttr>()))
+      AllowedInit = VD->getInit()->isConstantInitializer(
+          Context, VD->getType()->isReferenceType());
+
+    // Also make sure that destructor, if there is one, is empty.
+    if (AllowedInit)
+      if (CXXRecordDecl *RD = VD->getType()->getAsCXXRecordDecl())
+        AllowedInit =
+            isEmptyCudaDestructor(VD->getLocation(), RD->getDestructor());
+
+    if (!AllowedInit) {
+      Diag(VD->getLocation(), VD->hasAttr<CUDASharedAttr>()
+                                  ? diag::err_shared_var_init
+                                  : diag::err_dynamic_var_init)
+          << Init->getSourceRange();
+      VD->setInvalidDecl();
+    }
+  } else {
+    // This is a host-side global variable.  Check that the initializer is
+    // callable from the host side.
+    const FunctionDecl *InitFn = nullptr;
+    if (const CXXConstructExpr *CE = dyn_cast<CXXConstructExpr>(Init)) {
+      InitFn = CE->getConstructor();
+    } else if (const CallExpr *CE = dyn_cast<CallExpr>(Init)) {
+      InitFn = CE->getDirectCallee();
+    }
+    if (InitFn) {
+      CUDAFunctionTarget InitFnTarget = IdentifyCUDATarget(InitFn);
+      if (InitFnTarget != CFT_Host && InitFnTarget != CFT_HostDevice) {
+        Diag(VD->getLocation(), diag::err_ref_bad_target_global_initializer)
+            << InitFnTarget << InitFn;
+        Diag(InitFn->getLocation(), diag::note_previous_decl) << InitFn;
+        VD->setInvalidDecl();
+      }
+    }
+  }
+}
+
 // With -fcuda-host-device-constexpr, an unattributed constexpr function is
 // treated as implicitly __host__ __device__, unless:
 //  * it is a variadic function (device-side variadic functions are not
@@ -521,7 +576,7 @@ void Sema::maybeAddCUDAHostDeviceAttrs(FunctionDecl *NewD,
     if (!getSourceManager().isInSystemHeader(Match->getLocation())) {
       Diag(NewD->getLocation(),
            diag::err_cuda_unattributed_constexpr_cannot_overload_device)
-          << NewD->getName();
+          << NewD;
       Diag(Match->getLocation(),
            diag::note_cuda_conflicting_device_function_declared_here);
     }
@@ -790,9 +845,12 @@ bool Sema::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) {
   // If the caller is known-emitted, mark the callee as known-emitted.
   // Otherwise, mark the call in our call graph so we can traverse it later.
   bool CallerKnownEmitted = IsKnownEmitted(*this, Caller);
-  if (CallerKnownEmitted)
-    MarkKnownEmitted(*this, Caller, Callee, Loc);
-  else {
+  if (CallerKnownEmitted) {
+    // Host-side references to a __global__ function refer to the stub, so the
+    // function itself is never emitted and therefore should not be marked.
+    if (getLangOpts().CUDAIsDevice || IdentifyCUDATarget(Callee) != CFT_Global)
+      MarkKnownEmitted(*this, Caller, Callee, Loc);
+  } else {
     // If we have
     //   host fn calls kernel fn calls host+device,
     // the HD function does not get instantiated on the host.  We model this by