tra updated this revision to Diff 75462.
tra added a comment.
Addressed jlebar's comments.
https://reviews.llvm.org/D25796
Files:
lib/Sema/SemaExprCXX.cpp
test/SemaCUDA/overloaded-delete.cu
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;
}
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,
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
