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
