Index: cfe/trunk/lib/AST/RecordLayoutBuilder.cpp
===================================================================
--- cfe/trunk/lib/AST/RecordLayoutBuilder.cpp
+++ cfe/trunk/lib/AST/RecordLayoutBuilder.cpp
@@ -2025,6 +2025,21 @@
         continue;
     }
 
+    if (Context.getLangOpts().CUDA) {
+      // While compiler may see key method in this TU, during CUDA
+      // compilation we should ignore methods that are not accessible
+      // on this side of compilation.
+      if (Context.getLangOpts().CUDAIsDevice) {
+        // In device mode ignore methods without __device__ attribute.
+        if (!MD->hasAttr<CUDADeviceAttr>())
+          continue;
+      } else {
+        // In host mode ignore __device__-only methods.
+        if (!MD->hasAttr<CUDAHostAttr>() && MD->hasAttr<CUDADeviceAttr>())
+          continue;
+      }
+    }
+
     // If the key function is dllimport but the class isn't, then the class has
     // no key function. The DLL that exports the key function won't export the
     // vtable in this case.
Index: cfe/trunk/lib/CodeGen/CGVTables.cpp
===================================================================
--- cfe/trunk/lib/CodeGen/CGVTables.cpp
+++ cfe/trunk/lib/CodeGen/CGVTables.cpp
@@ -582,6 +582,24 @@
         break;
       }
 
+      if (CGM.getLangOpts().CUDA) {
+        // Emit NULL for methods we can't codegen on this
+        // side. Otherwise we'd end up with vtable with unresolved
+        // references.
+        const CXXMethodDecl *MD = cast<CXXMethodDecl>(GD.getDecl());
+        // OK on device side: functions w/ __device__ attribute
+        // OK on host side: anything except __device__-only functions.
+        bool CanEmitMethod = CGM.getLangOpts().CUDAIsDevice
+                                 ? MD->hasAttr<CUDADeviceAttr>()
+                                 : (MD->hasAttr<CUDAHostAttr>() ||
+                                    !MD->hasAttr<CUDADeviceAttr>());
+        if (!CanEmitMethod) {
+          Init = llvm::ConstantExpr::getNullValue(Int8PtrTy);
+          break;
+        }
+        // Method is acceptable, continue processing as usual.
+      }
+
       if (cast<CXXMethodDecl>(GD.getDecl())->isPure()) {
         // We have a pure virtual member function.
         if (!PureVirtualFn) {
Index: cfe/trunk/test/CodeGenCUDA/device-vtable.cu
===================================================================
--- cfe/trunk/test/CodeGenCUDA/device-vtable.cu
+++ cfe/trunk/test/CodeGenCUDA/device-vtable.cu
@@ -0,0 +1,61 @@
+// REQUIRES: x86-registered-target
+// REQUIRES: nvptx-registered-target
+
+// Make sure we don't emit vtables for classes with methods that have
+// inappropriate target attributes. Currently it's mostly needed in
+// order to avoid emitting vtables for host-only classes on device
+// side where we can't codegen them.
+
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm -o - %s \
+// RUN:     | FileCheck %s -check-prefix=CHECK-HOST -check-prefix=CHECK-BOTH
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -emit-llvm -o - %s \
+// RUN:     | FileCheck %s -check-prefix=CHECK-DEVICE -check-prefix=CHECK-BOTH
+
+#include "Inputs/cuda.h"
+
+struct H  {
+  virtual void method();
+};
+//CHECK-HOST: @_ZTV1H =
+//CHECK-HOST-SAME: @_ZN1H6methodEv
+//CHECK-DEVICE-NOT: @_ZTV1H =
+
+struct D  {
+   __device__ virtual void method();
+};
+
+//CHECK-DEVICE: @_ZTV1D
+//CHECK-DEVICE-SAME: @_ZN1D6methodEv
+//CHECK-HOST-NOT: @_ZTV1D
+
+// This is the case with mixed host and device virtual methods.  It's
+// impossible to emit a valid vtable in that case because only host or
+// only device methods would be available during host or device
+// compilation. At the moment Clang (and NVCC) emit NULL pointers for
+// unavailable methods,
+struct HD  {
+  virtual void h_method();
+  __device__ virtual void d_method();
+};
+// CHECK-BOTH: @_ZTV2HD
+// CHECK-DEVICE-NOT: @_ZN2HD8h_methodEv
+// CHECK-DEVICE-SAME: null
+// CHECK-DEVICE-SAME: @_ZN2HD8d_methodEv
+// CHECK-HOST-SAME: @_ZN2HD8h_methodEv
+// CHECK-HOST-NOT: @_ZN2HD8d_methodEv
+// CHECK-HOST-SAME: null
+// CHECK-BOTH-SAME: ]
+
+void H::method() {}
+//CHECK-HOST: define void @_ZN1H6methodEv
+
+void __device__ D::method() {}
+//CHECK-DEVICE: define void @_ZN1D6methodEv
+
+void __device__ HD::d_method() {}
+// CHECK-DEVICE: define void @_ZN2HD8d_methodEv
+// CHECK-HOST-NOT: define void @_ZN2HD8d_methodEv
+void HD::h_method() {}
+// CHECK-HOST: define void @_ZN2HD8h_methodEv
+// CHECK-DEVICE-NOT: define void @_ZN2HD8h_methodEv
+