Index: cfe/trunk/lib/Sema/SemaCUDA.cpp
===================================================================
--- cfe/trunk/lib/Sema/SemaCUDA.cpp
+++ cfe/trunk/lib/Sema/SemaCUDA.cpp
@@ -267,6 +267,18 @@
                                                    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 @@
     }
   }
 
+
+  // 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;
 }
Index: cfe/trunk/test/SemaCUDA/default-ctor.cu
===================================================================
--- cfe/trunk/test/SemaCUDA/default-ctor.cu
+++ cfe/trunk/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;
+}