r284879 - Declare H and H new/delete.

[Artem Belevich via cfe-commits]

Author: tra
Date: Fri Oct 21 15:34:05 2016
New Revision: 284879

URL: http://llvm.org/viewvc/llvm-project?rev=284879&view=rev
Log:
Declare H and H new/delete.

Modified:
    cfe/trunk/lib/Sema/SemaExprCXX.cpp
    cfe/trunk/test/SemaCUDA/overloaded-delete.cu

Modified: cfe/trunk/lib/Sema/SemaExprCXX.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaExprCXX.cpp?rev=284879&r1=284878&r2=284879&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaExprCXX.cpp (original)
+++ cfe/trunk/lib/Sema/SemaExprCXX.cpp Fri Oct 21 15:34:05 2016
@@ -2593,28 +2593,39 @@ void Sema::DeclareGlobalAllocationFuncti
         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();
-  }
-  Alloc->setParams(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.back()->setImplicit();
+    }
+    Alloc->setParams(ParamDecls);
+    if (ExtraAttr)
+      Alloc->addAttr(ExtraAttr);
+    Context.getTranslationUnitDecl()->addDecl(Alloc);
+    IdResolver.tryAddTopLevelDecl(Alloc, Name);
+  };
 
-  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,

Modified: cfe/trunk/test/SemaCUDA/overloaded-delete.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/overloaded-delete.cu?rev=284879&r1=284878&r2=284879&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/overloaded-delete.cu (original)
+++ cfe/trunk/test/SemaCUDA/overloaded-delete.cu Fri Oct 21 15:34:05 2016
@@ -16,10 +16,54 @@ __host__ __device__ void test(S* s) {
   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;
 }