[CUDA] Create __host__ and device variants of standard allocator declarations.

[tra]

Summary

Implicit functions are treated as if they were host device and clang does not allow overloading
those with host or device variants.

In order for users to provide their own standard allocators, we must create host and device variants of these declarations during CUDA compilation.

[jlebar]

Is it not possible to write a testcase for this?

[jlebar]

lib/Sema/SemaExprCXX.cpp

2596–2601

Nit, no need for empty parens.

2628–2629

I kind of feel like it would be easier to add a boolean parameter to CreateAllocationFunctionDecl, then you don't need to repeat these two lines.

[tra]

Addressed jlebar's comments.

[jlebar]

lib/Sema/SemaExprCXX.cpp

2596

Oh, I like this way better than a bool arg.

[tra]

Landed in r284879.

lib/Sema/SemaExprCXX.cpp

if (!getLangOpts().CPlusPlus11) {
       EPI.ExceptionSpec.Type = EST_Dynamic;
       EPI.ExceptionSpec.Exceptions = llvm::makeArrayRef(BadAllocType);
     }
   } else {
     EPI.ExceptionSpec =
         getLangOpts().CPlusPlus11 ? EST_BasicNoexcept : EST_DynamicNone;
   }
 
+  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);
+    FunctionDecl *Alloc = FunctionDecl::Create(
+        Context, GlobalCtx, SourceLocation(), SourceLocation(), Name,
     Alloc->setImplicit();
 
     // Implicit sized deallocation functions always have default visibility.
-  Alloc->addAttr(VisibilityAttr::CreateImplicit(Context,
-                                                VisibilityAttr::Default));
+    Alloc->addAttr(
+        VisibilityAttr::CreateImplicit(Context, VisibilityAttr::Default));
 
-  llvm::SmallVector<ParmVarDecl*, 3> ParamDecls;
+    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.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));
+  }
 }
 
 FunctionDecl *Sema::FindUsualDeallocationFunction(SourceLocation StartLoc,
                                                   bool CanProvideSize,
                                                   bool Overaligned,
                                                   DeclarationName Name) {
   DeclareGlobalNewDelete();
 
   LookupResult FoundDelete(*this, Name, StartLoc, LookupOrdinaryName);

test/SemaCUDA/overloaded-delete.cu

 };
 
 __host__ __device__ void test(S* s) {
   // This shouldn't be ambiguous -- we call the host overload in host mode and
   // the device overload in device mode.
   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;
 }