[CUDA][HIP] Fix hostness of defaulted constructor

Clang does not respect the explicit device host attributes of defaulted special members.
Also clang does not respect the hostness of special members determined by their
first declarations.
Clang also adds duplicate implicit device or host attributes in certain cases.
This patch fixes that.
Differential Revision: https://reviews.llvm.org/D67509


git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@372394 91177308-0d34-0410-b5e6-96231b3b80d8

2 files changed, 70 insertions, 12 deletions

diff --git a/lib/Sema/SemaCUDA.cpp b/lib/Sema/SemaCUDA.cpp
index 203c09c571..e009dcb6f1 100644
--- a/lib/Sema/SemaCUDA.cpp
+++ b/lib/Sema/SemaCUDA.cpp
@@ -267,6 +267,18 @@ bool Sema::inferCUDATargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl,
                                                    CXXMethodDecl *MemberDecl,
                                                    bool ConstRHS,
                                                    bool Diagnose) {
+  // If the defaulted special member is defined lexically outside of its
+  // owning class, or the special member already has explicit device or host
+  // attributes, do not infer.
+  bool InClass = MemberDecl->getLexicalParent() == MemberDecl->getParent();
+  bool HasH = MemberDecl->hasAttr<CUDAHostAttr>();
+  bool HasD = MemberDecl->hasAttr<CUDADeviceAttr>();
+  bool HasExplicitAttr =
+      (HasD && !MemberDecl->getAttr<CUDADeviceAttr>()->isImplicit()) ||
+      (HasH && !MemberDecl->getAttr<CUDAHostAttr>()->isImplicit());
+  if (!InClass || HasExplicitAttr)
+    return false;
+
   llvm::Optional<CUDAFunctionTarget> InferredTarget;
 
   // We're going to invoke special member lookup; mark that these special
@@ -371,21 +383,24 @@ bool Sema::inferCUDATargetForImplicitSpecialMember(CXXRecordDecl *ClassDecl,
     }
   }
 
+
+  // If no target was inferred, mark this member as __host__ __device__;
+  // it's the least restrictive option that can be invoked from any target.
+  bool NeedsH = true, NeedsD = true;
   if (InferredTarget.hasValue()) {
-    if (InferredTarget.getValue() == CFT_Device) {
-      MemberDecl->addAttr(CUDADeviceAttr::CreateImplicit(Context));
-    } else if (InferredTarget.getValue() == CFT_Host) {
-      MemberDecl->addAttr(CUDAHostAttr::CreateImplicit(Context));
-    } else {
-      MemberDecl->addAttr(CUDADeviceAttr::CreateImplicit(Context));
-      MemberDecl->addAttr(CUDAHostAttr::CreateImplicit(Context));
-    }
-  } else {
-    // If no target was inferred, mark this member as __host__ __device__;
-    // it's the least restrictive option that can be invoked from any target.
+    if (InferredTarget.getValue() == CFT_Device)
+      NeedsH = false;
+    else if (InferredTarget.getValue() == CFT_Host)
+      NeedsD = false;
+  }
+
+  // We either setting attributes first time, or the inferred ones must match
+  // previously set ones.
+  assert(!(HasD || HasH) || (NeedsD == HasD && NeedsH == HasH));
+  if (NeedsD && !HasD)
     MemberDecl->addAttr(CUDADeviceAttr::CreateImplicit(Context));
+  if (NeedsH && !HasH)
     MemberDecl->addAttr(CUDAHostAttr::CreateImplicit(Context));
-  }
 
   return false;
 }
diff --git a/test/SemaCUDA/default-ctor.cu b/test/SemaCUDA/default-ctor.cu
new file mode 100644
index 0000000000..cbad7a1774
--- /dev/null
+++ b/test/SemaCUDA/default-ctor.cu
@@ -0,0 +1,43 @@
+// RUN: %clang_cc1 -std=c++11 -triple nvptx64-nvidia-cuda -fsyntax-only \
+// RUN:            -fcuda-is-device -verify -verify-ignore-unexpected=note %s
+// RUN: %clang_cc1 -std=c++11 -triple x86_64-unknown-linux-gnu -fsyntax-only \
+// RUN:            -verify -verify-ignore-unexpected=note %s
+
+#include "Inputs/cuda.h"
+
+struct In { In() = default; };
+struct InD { __device__ InD() = default; };
+struct InH { __host__ InH() = default; };
+struct InHD { __host__ __device__ InHD() = default; };
+
+struct Out { Out(); };
+struct OutD { __device__ OutD(); };
+struct OutH { __host__ OutH(); };
+struct OutHD { __host__ __device__ OutHD(); };
+
+Out::Out() = default;
+__device__ OutD::OutD() = default;
+__host__ OutH::OutH() = default;
+__host__ __device__ OutHD::OutHD() = default;
+
+__device__ void fd() {
+  In in;
+  InD ind;
+  InH inh; // expected-error{{no matching constructor for initialization of 'InH'}}
+  InHD inhd;
+  Out out; // expected-error{{no matching constructor for initialization of 'Out'}}
+  OutD outd;
+  OutH outh; // expected-error{{no matching constructor for initialization of 'OutH'}}
+  OutHD outhd;
+}
+
+__host__ void fh() {
+  In in;
+  InD ind; // expected-error{{no matching constructor for initialization of 'InD'}}
+  InH inh;
+  InHD inhd;
+  Out out;
+  OutD outd; // expected-error{{no matching constructor for initialization of 'OutD'}}
+  OutH outh;
+  OutHD outhd;
+}