Index: clang/lib/Sema/SemaDecl.cpp
===================================================================
--- clang/lib/Sema/SemaDecl.cpp
+++ clang/lib/Sema/SemaDecl.cpp
@@ -17746,6 +17746,17 @@
   if (FD->isDependentContext())
     return FunctionEmissionStatus::TemplateDiscarded;
 
+  // The Microsoft ABI requires that we perform the destructor body
+  // checks (i.e. operator delete() lookup) when the vtable is marked used, as
+  // the deleting destructor is emitted with the vtable. Such check may happen
+  // even though the destructor is not defined yet.
+  if (Context.getTargetInfo().getCXXABI().isMicrosoft()) {
+    if (auto *DD = dyn_cast<CXXDestructorDecl>(FD)) {
+      if (DD == CurContext && !DD->getDefinition())
+        return FunctionEmissionStatus::Emitted;
+    }
+  }
+
   FunctionEmissionStatus OMPES = FunctionEmissionStatus::Unknown;
   if (LangOpts.OpenMPIsDevice) {
     Optional<OMPDeclareTargetDeclAttr::DevTypeTy> DevTy =
Index: clang/test/SemaCUDA/deleting-dtor.cu
===================================================================
--- /dev/null
+++ clang/test/SemaCUDA/deleting-dtor.cu
@@ -0,0 +1,91 @@
+// RUN: %clang_cc1 -triple x86_64-pc-windows-msvc -fsyntax-only -verify=MS -verify=GEN %s
+// RUN: %clang_cc1 -triple x86_64-pc-linux -fsyntax-only -verify=GEN %s
+
+#include "Inputs/cuda.h"
+typedef long long __m256i __attribute__((__vector_size__(32)));
+
+namespace NoDiag {
+  struct CFileStream {
+    CFileStream();
+    virtual ~CFileStream();
+  };
+
+  struct CMultiFileStream {
+    CFileStream m_fileStream;
+    ~CMultiFileStream();
+  };
+
+  CFileStream::CFileStream() {}
+  CFileStream::~CFileStream() {}
+  CMultiFileStream::~CMultiFileStream() {}
+}
+
+// No diagnostic since deleting dtor is not emitted.
+namespace NoVtbl {
+// Intentionally generates delayed diagnostic about r0.
+// This diagnostic is not supposed to be emitted unless f is emitted.
+  static __device__ __host__ void f(__m256i *p) {
+    __asm__ volatile("vmovaps  %0, %%ymm0" ::"m"(*(__m256i *)p)
+                 : "r0");
+  }
+  struct CFileStream {
+    void operator delete(void *p) {
+      f(0);
+    }
+    CFileStream();
+    virtual ~CFileStream();
+  };
+
+  struct CMultiFileStream {
+    CFileStream m_fileStream;
+    ~CMultiFileStream();
+  };
+}
+
+// Only MS has diagnostic since MS requires deleting dtor to be emitted when
+// vtable is emitted, even though dtor is not defined.
+namespace MSEmitDeletingDtor {
+  static __device__ __host__ void f(__m256i *p) {
+    __asm__ volatile("vmovaps  %0, %%ymm0" ::"m"(*(__m256i *)p)
+                   : "r0"); // MS-error{{unknown register name 'r0' in asm}}
+  }
+  struct CFileStream {
+    void operator delete(void *p) {
+      f(0);  // MS-note{{called by 'operator delete'}}
+    }
+    CFileStream();
+    virtual ~CFileStream();  // MS-note{{called by '~CFileStream'}}
+  };
+
+  struct CMultiFileStream {
+    CFileStream m_fileStream;
+    ~CMultiFileStream();
+  };
+
+  // This causes vtable emitted so that deleting dtor is emitted for MS.
+  CFileStream::CFileStream() {}
+}
+
+namespace EmitDtor {
+  static __device__ __host__ void f(__m256i *p) {
+    __asm__ volatile("vmovaps  %0, %%ymm0" ::"m"(*(__m256i *)p)
+                 : "r0"); // GEN-error{{unknown register name 'r0' in asm}}
+  }
+  struct CFileStream {
+    void operator delete(void *p) {
+      f(0);
+    }
+    CFileStream();
+    virtual ~CFileStream();
+  };
+
+  struct CMultiFileStream {
+    CFileStream m_fileStream;
+    ~CMultiFileStream();
+  };
+
+  CFileStream::~CFileStream() {
+    f(0); // GEN-note{{called by '~CFileStream'}}
+  }
+}
+