Index: lib/AST/RecordLayoutBuilder.cpp
===================================================================
--- lib/AST/RecordLayoutBuilder.cpp
+++ lib/AST/RecordLayoutBuilder.cpp
@@ -1996,6 +1996,16 @@
   bool allowInlineFunctions =
     Context.getTargetInfo().getCXXABI().canKeyFunctionBeInline();
 
+  if (Context.getLangOpts().CUDA) {
+    const bool IsDevice = Context.getLangOpts().CUDAIsDevice;
+    for (const CXXMethodDecl *MD : RD->methods())
+      if (Context.getLangOpts().CUDAIsDevice && !MD->hasAttr<CUDADeviceAttr>())
+        return nullptr;
+      else if (!Context.getLangOpts().CUDAIsDevice &&
+               !MD->hasAttr<CUDAHostAttr>() && MD->hasAttr<CUDADeviceAttr>())
+        return nullptr;
+  }
+
   for (const CXXMethodDecl *MD : RD->methods()) {
     if (!MD->isVirtual())
       continue;
Index: test/CodeGenCUDA/device-vtable.cu
===================================================================
--- /dev/null
+++ test/CodeGenCUDA/device-vtable.cu
@@ -0,0 +1,55 @@
+// 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
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -emit-llvm -o - %s \
+// RUN:     | FileCheck %s -check-prefix=CHECK-DEVICE
+
+#include "Inputs/cuda.h"
+
+class H  {
+ public:
+  virtual void method();
+};
+//CHECK-HOST: @_ZTV1H =
+//CHECK-HOST-SAME: @_ZN1H6methodEv
+//CHECK-DEVICE-NOT: @_ZTV1H =
+
+class D  {
+ public:
+   __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. For now we'll not emit such vtable at all.
+class HD  {
+ public:
+   virtual void h_method();
+   __device__ virtual void d_method();
+};
+//CHECK-BOTH-NOT: @_ZTV2HD
+
+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