r372394 - [CUDA][HIP] Fix hostness of defaulted constructor

[Yaxun Liu via cfe-commits]

Author: yaxunl
Date: Fri Sep 20 07:28:09 2019
New Revision: 372394

URL: http://llvm.org/viewvc/llvm-project?rev=372394&view=rev
Log:
[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

Added:
    cfe/trunk/test/SemaCUDA/default-ctor.cu
Modified:
    cfe/trunk/lib/Sema/SemaCUDA.cpp

Modified: cfe/trunk/lib/Sema/SemaCUDA.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaCUDA.cpp?rev=372394&r1=372393&r2=372394&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaCUDA.cpp (original)
+++ cfe/trunk/lib/Sema/SemaCUDA.cpp Fri Sep 20 07:28:09 2019
@@ -267,6 +267,18 @@ bool Sema::inferCUDATargetForImplicitSpe
                                                    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::inferCUDATargetForImplicitSpe
     }
   }
 
+
+  // 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;
 }

Added: cfe/trunk/test/SemaCUDA/default-ctor.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/default-ctor.cu?rev=372394&view=auto
==============================================================================
--- cfe/trunk/test/SemaCUDA/default-ctor.cu (added)
+++ cfe/trunk/test/SemaCUDA/default-ctor.cu Fri Sep 20 07:28:09 2019
@@ -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;
+}