Index: lib/Sema/SemaOverload.cpp
===================================================================
--- lib/Sema/SemaOverload.cpp
+++ lib/Sema/SemaOverload.cpp
@@ -5977,6 +5977,15 @@
     }
   }
 
+  // (CUDA B.1): Check for invalid calls between targets.
+  if (getLangOpts().CUDA)
+    if (const FunctionDecl *Caller = dyn_cast<FunctionDecl>(CurContext))
+      if (CheckCUDATarget(Caller, Method)) {
+        Candidate.Viable = false;
+        Candidate.FailureKind = ovl_fail_bad_target;
+        return;
+      }
+
   // Determine the implicit conversion sequences for each of the
   // arguments.
   for (unsigned ArgIdx = 0; ArgIdx < Args.size(); ++ArgIdx) {
@@ -11595,6 +11604,18 @@
     new (Context) CXXMemberCallExpr(Context, MemExprE, Args,
                                     ResultType, VK, RParenLoc);
 
+  // (CUDA B.1): Check for invalid calls between targets.
+  if (getLangOpts().CUDA) {
+    if (const FunctionDecl *Caller = dyn_cast<FunctionDecl>(CurContext)) {
+      if (CheckCUDATarget(Caller, Method)) {
+        Diag(MemExpr->getMemberLoc(), diag::err_ref_bad_target)
+            << IdentifyCUDATarget(Method) << Method->getIdentifier()
+            << IdentifyCUDATarget(Caller);
+        return ExprError();
+      }
+    }
+  }
+
   // Check for a valid return type.
   if (CheckCallReturnType(Method->getReturnType(), MemExpr->getMemberLoc(),
                           TheCall, Method))
Index: test/SemaCUDA/method-target.cu
===================================================================
--- /dev/null
+++ test/SemaCUDA/method-target.cu
@@ -0,0 +1,71 @@
+// RUN: %clang_cc1 -fsyntax-only -verify %s
+
+#include "Inputs/cuda.h"
+
+//------------------------------------------------------------------------------
+// Test 1: host method called from device function
+
+struct S1 {
+  void method() {}
+};
+
+__device__ void foo1(S1& s) {
+  s.method(); // expected-error {{reference to __host__ function 'method' in __device__ function}}
+}
+
+//------------------------------------------------------------------------------
+// Test 2: host method called from device function, for overloaded method
+
+struct S2 {
+  void method(int) {} // expected-note {{candidate function not viable: call to __host__ function from __device__ function}}
+  void method(float) {} // expected-note {{candidate function not viable: call to __host__ function from __device__ function}}
+};
+
+__device__ void foo2(S2& s, int i, float f) {
+  s.method(f); // expected-error {{no matching member function}}
+}
+
+//------------------------------------------------------------------------------
+// Test 3: device method called from host function
+
+struct S3 {
+  __device__ void method() {}
+};
+
+void foo3(S3& s) {
+  s.method(); // expected-error {{reference to __device__ function 'method' in __host__ function}}
+}
+
+//------------------------------------------------------------------------------
+// Test 4: device method called from host&device function
+
+struct S4 {
+  __device__ void method() {}
+};
+
+__host__ __device__ void foo4(S4& s) {
+  s.method(); // expected-error {{reference to __device__ function 'method' in __host__ __device__ function}}
+}
+
+//------------------------------------------------------------------------------
+// Test 5: overloaded operators
+
+struct S5 {
+  S5() {}
+  S5& operator=(const S5&) {return *this;} // expected-note {{candidate function not viable}}
+};
+
+__device__ void foo5(S5& s, S5& t) {
+  s = t; // expected-error {{no viable overloaded '='}}
+}
+
+//------------------------------------------------------------------------------
+// Test 6: call method through pointer
+
+struct S6 {
+  void method() {}
+};
+
+__device__ void foo6(S6* s) {
+  s->method(); // expected-error {{reference to __host__ function 'method' in __device__ function}}
+}