Index: cfe/trunk/lib/Sema/SemaDeclCXX.cpp
===================================================================
--- cfe/trunk/lib/Sema/SemaDeclCXX.cpp
+++ cfe/trunk/lib/Sema/SemaDeclCXX.cpp
@@ -7222,8 +7222,17 @@
   if (getLangOpts().CUDA) {
     // We should delete the special member in CUDA mode if target inference
     // failed.
-    return inferCUDATargetForImplicitSpecialMember(RD, CSM, MD, SMI.ConstArg,
-                                                   Diagnose);
+    // For inherited constructors (non-null ICI), CSM may be passed so that MD
+    // is treated as certain special member, which may not reflect what special
+    // member MD really is. However inferCUDATargetForImplicitSpecialMember
+    // expects CSM to match MD, therefore recalculate CSM.
+    assert(ICI || CSM == getSpecialMember(MD));
+    auto RealCSM = CSM;
+    if (ICI)
+      RealCSM = getSpecialMember(MD);
+
+    return inferCUDATargetForImplicitSpecialMember(RD, RealCSM, MD,
+                                                   SMI.ConstArg, Diagnose);
   }
 
   return false;
Index: cfe/trunk/test/SemaCUDA/implicit-member-target-inherited.cu
===================================================================
--- cfe/trunk/test/SemaCUDA/implicit-member-target-inherited.cu
+++ cfe/trunk/test/SemaCUDA/implicit-member-target-inherited.cu
@@ -0,0 +1,205 @@
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify %s -Wno-defaulted-function-deleted
+
+#include "Inputs/cuda.h"
+
+//------------------------------------------------------------------------------
+// Test 1: infer inherited default ctor to be host.
+
+struct A1_with_host_ctor {
+  A1_with_host_ctor() {}
+};
+// expected-note@-3 {{candidate constructor (the implicit copy constructor) not viable}}
+// expected-note@-4 {{candidate constructor (the implicit move constructor) not viable}}
+
+// The inherited default constructor is inferred to be host, so we'll encounter
+// an error when calling it from a __device__ function, but not from a __host__
+// function.
+struct B1_with_implicit_default_ctor : A1_with_host_ctor {
+  using A1_with_host_ctor::A1_with_host_ctor;
+};
+
+// expected-note@-4 {{call to __host__ function from __device__}}
+// expected-note@-5 {{candidate constructor (the implicit copy constructor) not viable}}
+// expected-note@-6 {{candidate constructor (the implicit move constructor) not viable}}
+// expected-note@-6 2{{constructor from base class 'A1_with_host_ctor' inherited here}}
+
+void hostfoo() {
+  B1_with_implicit_default_ctor b;
+}
+
+__device__ void devicefoo() {
+  B1_with_implicit_default_ctor b; // expected-error {{no matching constructor}}
+}
+
+//------------------------------------------------------------------------------
+// Test 2: infer inherited default ctor to be device.
+
+struct A2_with_device_ctor {
+  __device__ A2_with_device_ctor() {}
+};
+// expected-note@-3 {{candidate constructor (the implicit copy constructor) not viable}}
+// expected-note@-4 {{candidate constructor (the implicit move constructor) not viable}}
+
+struct B2_with_implicit_default_ctor : A2_with_device_ctor {
+  using A2_with_device_ctor::A2_with_device_ctor;
+};
+
+// expected-note@-4 {{call to __device__ function from __host__}}
+// expected-note@-5 {{candidate constructor (the implicit copy constructor) not viable}}
+// expected-note@-6 {{candidate constructor (the implicit move constructor) not viable}}
+// expected-note@-6 2{{constructor from base class 'A2_with_device_ctor' inherited here}}
+
+void hostfoo2() {
+  B2_with_implicit_default_ctor b;  // expected-error {{no matching constructor}}
+}
+
+__device__ void devicefoo2() {
+  B2_with_implicit_default_ctor b;
+}
+
+//------------------------------------------------------------------------------
+// Test 3: infer inherited copy ctor
+
+struct A3_with_device_ctors {
+  __host__ A3_with_device_ctors() {}
+  __device__ A3_with_device_ctors(const A3_with_device_ctors&) {}
+};
+
+struct B3_with_implicit_ctors : A3_with_device_ctors {
+  using A3_with_device_ctors::A3_with_device_ctors;
+};
+// expected-note@-3 2{{call to __device__ function from __host__ function}}
+// expected-note@-4 {{default constructor}}
+
+
+void hostfoo3() {
+  B3_with_implicit_ctors b;  // this is OK because the inferred inherited default ctor
+                             // here is __host__
+  B3_with_implicit_ctors b2 = b; // expected-error {{no matching constructor}}
+
+}
+
+//------------------------------------------------------------------------------
+// Test 4: infer inherited default ctor from a field, not a base
+
+struct A4_with_host_ctor {
+  A4_with_host_ctor() {}
+};
+
+struct B4_with_inherited_host_ctor : A4_with_host_ctor{
+  using A4_with_host_ctor::A4_with_host_ctor;
+};
+
+struct C4_with_implicit_default_ctor {
+  B4_with_inherited_host_ctor field;
+};
+
+// expected-note@-4 {{call to __host__ function from __device__}}
+// expected-note@-5 {{candidate constructor (the implicit copy constructor) not viable}}
+// expected-note@-6 {{candidate constructor (the implicit move constructor) not viable}}
+
+void hostfoo4() {
+  C4_with_implicit_default_ctor b;
+}
+
+__device__ void devicefoo4() {
+  C4_with_implicit_default_ctor b; // expected-error {{no matching constructor}}
+}
+
+//------------------------------------------------------------------------------
+// Test 5: inherited copy ctor with non-const param
+
+struct A5_copy_ctor_constness {
+  __host__ A5_copy_ctor_constness() {}
+  __host__ A5_copy_ctor_constness(A5_copy_ctor_constness&) {}
+};
+
+struct B5_copy_ctor_constness : A5_copy_ctor_constness {
+  using A5_copy_ctor_constness::A5_copy_ctor_constness;
+};
+
+// expected-note@-4 {{candidate constructor (the implicit copy constructor) not viable: call to __host__ function from __device__ function}}
+// expected-note@-5 {{candidate constructor (the implicit default constructor) not viable}}
+
+void hostfoo5(B5_copy_ctor_constness& b_arg) {
+  B5_copy_ctor_constness b = b_arg;
+}
+
+__device__ void devicefoo5(B5_copy_ctor_constness& b_arg) {
+  B5_copy_ctor_constness b = b_arg; // expected-error {{no matching constructor}}
+}
+
+//------------------------------------------------------------------------------
+// Test 6: explicitly defaulted ctor
+
+struct A6_with_device_ctor {
+  __device__ A6_with_device_ctor() {}
+};
+
+struct B6_with_defaulted_ctor : A6_with_device_ctor {
+  using A6_with_device_ctor::A6_with_device_ctor;
+  __host__ B6_with_defaulted_ctor() = default;
+};
+
+// expected-note@-3 {{explicitly defaulted function was implicitly deleted here}}
+// expected-note@-6 {{default constructor of 'B6_with_defaulted_ctor' is implicitly deleted because base class 'A6_with_device_ctor' has no default constructor}}
+
+void hostfoo6() {
+  B6_with_defaulted_ctor b; // expected-error {{call to implicitly-deleted default constructor}}
+}
+
+__device__ void devicefoo6() {
+  B6_with_defaulted_ctor b;
+}
+
+//------------------------------------------------------------------------------
+// Test 7: inherited copy assignment operator
+
+struct A7_with_copy_assign {
+  A7_with_copy_assign() {}
+  __device__ A7_with_copy_assign& operator=(const A7_with_copy_assign&) {}
+};
+
+struct B7_with_copy_assign : A7_with_copy_assign {
+  using A7_with_copy_assign::A7_with_copy_assign;
+};
+
+// expected-note@-4 {{candidate function (the implicit copy assignment operator) not viable: call to __device__ function from __host__ function}}
+// expected-note@-5 {{candidate function (the implicit move assignment operator) not viable: call to __device__ function from __host__ function}}
+
+void hostfoo7() {
+  B7_with_copy_assign b1, b2;
+  b1 = b2; // expected-error {{no viable overloaded '='}}
+}
+
+//------------------------------------------------------------------------------
+// Test 8: inherited move assignment operator
+
+// definitions for std::move
+namespace std {
+inline namespace foo {
+template <class T> struct remove_reference { typedef T type; };
+template <class T> struct remove_reference<T&> { typedef T type; };
+template <class T> struct remove_reference<T&&> { typedef T type; };
+
+template <class T> typename remove_reference<T>::type&& move(T&& t);
+}
+}
+
+struct A8_with_move_assign {
+  A8_with_move_assign() {}
+  __device__ A8_with_move_assign& operator=(A8_with_move_assign&&) {}
+  __device__ A8_with_move_assign& operator=(const A8_with_move_assign&) {}
+};
+
+struct B8_with_move_assign : A8_with_move_assign {
+  using A8_with_move_assign::A8_with_move_assign;
+};
+
+// expected-note@-4 {{candidate function (the implicit copy assignment operator) not viable: call to __device__ function from __host__ function}}
+// expected-note@-5 {{candidate function (the implicit move assignment operator) not viable: call to __device__ function from __host__ function}}
+
+void hostfoo8() {
+  B8_with_move_assign b1, b2;
+  b1 = std::move(b2); // expected-error {{no viable overloaded '='}}
+}
Index: cfe/trunk/test/SemaCUDA/inherited-ctor.cu
===================================================================
--- cfe/trunk/test/SemaCUDA/inherited-ctor.cu
+++ cfe/trunk/test/SemaCUDA/inherited-ctor.cu
@@ -0,0 +1,89 @@
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify %s
+
+// Inherit a valid non-default ctor.
+namespace NonDefaultCtorValid {
+  struct A {
+    A(const int &x) {}
+  };
+
+  struct B : A {
+    using A::A;
+  };
+
+  struct C {
+    struct B b;
+    C() : b(0) {}
+  };
+
+  void test() {
+    B b(0);
+    C c;
+  }
+}
+
+// Inherit an invalid non-default ctor.
+// The inherited ctor is invalid because it is unable to initialize s.
+namespace NonDefaultCtorInvalid {
+  struct S {
+    S() = delete;
+  };
+  struct A {
+    A(const int &x) {}
+  };
+
+  struct B : A {
+    using A::A;
+    S s;
+  };
+
+  struct C {
+    struct B b;
+    C() : b(0) {} // expected-error{{constructor inherited by 'B' from base class 'A' is implicitly deleted}}
+                  // expected-note@-6{{constructor inherited by 'B' is implicitly deleted because field 's' has a deleted corresponding constructor}}
+                  // expected-note@-15{{'S' has been explicitly marked deleted here}}
+  };
+}
+
+// Inherit a valid default ctor.
+namespace DefaultCtorValid {
+  struct A {
+    A() {}
+  };
+
+  struct B : A {
+    using A::A;
+  };
+
+  struct C {
+    struct B b;
+    C() {}
+  };
+
+  void test() {
+    B b;
+    C c;
+  }
+}
+
+// Inherit an invalid default ctor.
+// The inherited ctor is invalid because it is unable to initialize s.
+namespace DefaultCtorInvalid {
+  struct S {
+    S() = delete;
+  };
+  struct A {
+    A() {}
+  };
+
+  struct B : A {
+    using A::A;
+    S s;
+  };
+
+  struct C {
+    struct B b;
+    C() {} // expected-error{{call to implicitly-deleted default constructor of 'struct B'}}
+           // expected-note@-6{{default constructor of 'B' is implicitly deleted because field 's' has a deleted default constructor}}
+           // expected-note@-15{{'S' has been explicitly marked deleted here}}
+  };
+}