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; } _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits