Index: lib/Sema/SemaExprCXX.cpp
===================================================================
--- lib/Sema/SemaExprCXX.cpp
+++ lib/Sema/SemaExprCXX.cpp
@@ -2593,28 +2593,39 @@
         getLangOpts().CPlusPlus11 ? EST_BasicNoexcept : EST_DynamicNone;
   }
 
-  QualType FnType = Context.getFunctionType(Return, Params, EPI);
-  FunctionDecl *Alloc =
-    FunctionDecl::Create(Context, GlobalCtx, SourceLocation(),
-                         SourceLocation(), Name,
-                         FnType, /*TInfo=*/nullptr, SC_None, false, true);
-  Alloc->setImplicit();
+  auto CreateAllocationFunctionDecl = [&](Attr *ExtraAttr) {
+    QualType FnType = Context.getFunctionType(Return, Params, EPI);
+    FunctionDecl *Alloc = FunctionDecl::Create(
+        Context, GlobalCtx, SourceLocation(), SourceLocation(), Name,
+        FnType, /*TInfo=*/nullptr, SC_None, false, true);
+    Alloc->setImplicit();
 
-  // Implicit sized deallocation functions always have default visibility.
-  Alloc->addAttr(VisibilityAttr::CreateImplicit(Context,
-                                                VisibilityAttr::Default));
+    // Implicit sized deallocation functions always have default visibility.
+    Alloc->addAttr(
+        VisibilityAttr::CreateImplicit(Context, VisibilityAttr::Default));
 
-  llvm::SmallVector<ParmVarDecl*, 3> ParamDecls;
-  for (QualType T : Params) {
-    ParamDecls.push_back(
-        ParmVarDecl::Create(Context, Alloc, SourceLocation(), SourceLocation(),
-                            nullptr, T, /*TInfo=*/nullptr, SC_None, nullptr));
-    ParamDecls.back()->setImplicit();
+    llvm::SmallVector<ParmVarDecl *, 3> ParamDecls;
+    for (QualType T : Params) {
+      ParamDecls.push_back(ParmVarDecl::Create(
+          Context, Alloc, SourceLocation(), SourceLocation(), nullptr, T,
+          /*TInfo=*/nullptr, SC_None, nullptr));
+      ParamDecls.back()->setImplicit();
+    }
+    Alloc->setParams(ParamDecls);
+    if (ExtraAttr)
+      Alloc->addAttr(ExtraAttr);
+    Context.getTranslationUnitDecl()->addDecl(Alloc);
+    IdResolver.tryAddTopLevelDecl(Alloc, Name);
+  };
+
+  if (!LangOpts.CUDA)
+    CreateAllocationFunctionDecl(nullptr);
+  else {
+    // Host and device get their own declaration so each can be
+    // defined or re-declared independently.
+    CreateAllocationFunctionDecl(CUDAHostAttr::CreateImplicit(Context));
+    CreateAllocationFunctionDecl(CUDADeviceAttr::CreateImplicit(Context));
   }
-  Alloc->setParams(ParamDecls);
-
-  Context.getTranslationUnitDecl()->addDecl(Alloc);
-  IdResolver.tryAddTopLevelDecl(Alloc, Name);
 }
 
 FunctionDecl *Sema::FindUsualDeallocationFunction(SourceLocation StartLoc,
Index: test/SemaCUDA/overloaded-delete.cu
===================================================================
--- test/SemaCUDA/overloaded-delete.cu
+++ test/SemaCUDA/overloaded-delete.cu
@@ -16,10 +16,54 @@
   delete s;
 }
 
+// Code should work with no explicit declarations/definitions of
+// allocator functions.
+__host__ __device__ void test_default_global_delete_hd(int *ptr) {
+  // Again, there should be no ambiguity between which operator delete we call.
+  ::delete ptr;
+}
+
+__device__ void test_default_global_delete(int *ptr) {
+  // Again, there should be no ambiguity between which operator delete we call.
+  ::delete ptr;
+}
+__host__ void test_default_global_delete(int *ptr) {
+  // Again, there should be no ambiguity between which operator delete we call.
+  ::delete ptr;
+}
+
+// It should work with only some of allocators (re-)declared.
+__device__ void operator delete(void *ptr);
+
+__host__ __device__ void test_partial_global_delete_hd(int *ptr) {
+  // Again, there should be no ambiguity between which operator delete we call.
+  ::delete ptr;
+}
+
+__device__ void test_partial_global_delete(int *ptr) {
+  // Again, there should be no ambiguity between which operator delete we call.
+  ::delete ptr;
+}
+__host__ void test_partial_global_delete(int *ptr) {
+  // Again, there should be no ambiguity between which operator delete we call.
+  ::delete ptr;
+}
+
+
+// We should be able to define both host and device variants.
 __host__ void operator delete(void *ptr) {}
 __device__ void operator delete(void *ptr) {}
 
-__host__ __device__ void test_global_delete(int *ptr) {
+__host__ __device__ void test_overloaded_global_delete_hd(int *ptr) {
+  // Again, there should be no ambiguity between which operator delete we call.
+  ::delete ptr;
+}
+
+__device__ void test_overloaded_global_delete(int *ptr) {
+  // Again, there should be no ambiguity between which operator delete we call.
+  ::delete ptr;
+}
+__host__ void test_overloaded_global_delete(int *ptr) {
   // Again, there should be no ambiguity between which operator delete we call.
   ::delete ptr;
 }