From 78929efb4dcbb2c8aeba46eee912a614d4ecebb9 Mon Sep 17 00:00:00 2001 From: Artem Belevich Date: Fri, 21 Sep 2018 17:29:33 +0000 Subject: [PATCH] [CUDA] Ignore uncallable functions when we check for usual deallocators. Previously clang considered function variants from both sides of compilation and that resulted in picking up wrong deallocation function. Differential Revision: https://reviews.llvm.org/D51808 llvm-svn: 342749 --- clang/include/clang/AST/DeclCXX.h | 13 ++- clang/include/clang/Sema/Sema.h | 2 + clang/lib/AST/DeclCXX.cpp | 20 ++-- clang/lib/Sema/SemaDeclCXX.cpp | 3 +- clang/lib/Sema/SemaExprCXX.cpp | 24 ++++- clang/test/CodeGenCUDA/usual-deallocators.cu | 133 ++++++++++++++++++++++++ clang/test/SemaCUDA/call-host-fn-from-device.cu | 6 +- clang/test/SemaCUDA/usual-deallocators.cu | 95 +++++++++++++++++ 8 files changed, 279 insertions(+), 17 deletions(-) create mode 100644 clang/test/CodeGenCUDA/usual-deallocators.cu create mode 100644 clang/test/SemaCUDA/usual-deallocators.cu diff --git a/clang/include/clang/AST/DeclCXX.h b/clang/include/clang/AST/DeclCXX.h index e0c6c00..d531ee1 100644 --- a/clang/include/clang/AST/DeclCXX.h +++ b/clang/include/clang/AST/DeclCXX.h @@ -2109,10 +2109,15 @@ public: Base, IsAppleKext); } - /// Determine whether this is a usual deallocation function - /// (C++ [basic.stc.dynamic.deallocation]p2), which is an overloaded - /// delete or delete[] operator with a particular signature. - bool isUsualDeallocationFunction() const; + /// Determine whether this is a usual deallocation function (C++ + /// [basic.stc.dynamic.deallocation]p2), which is an overloaded delete or + /// delete[] operator with a particular signature. Populates \p PreventedBy + /// with the declarations of the functions of the same kind if they were the + /// reason for this function returning false. This is used by + /// Sema::isUsualDeallocationFunction to reconsider the answer based on the + /// context. + bool isUsualDeallocationFunction( + SmallVectorImpl &PreventedBy) const; /// Determine whether this is a copy-assignment operator, regardless /// of whether it was declared implicitly or explicitly. diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index a701876..f91331b 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -1619,6 +1619,8 @@ public: SourceLocation Loc, const NamedDecl *D, ArrayRef Equiv); + bool isUsualDeallocationFunction(const CXXMethodDecl *FD); + bool isCompleteType(SourceLocation Loc, QualType T) { return !RequireCompleteTypeImpl(Loc, T, nullptr); } diff --git a/clang/lib/AST/DeclCXX.cpp b/clang/lib/AST/DeclCXX.cpp index 8582b54..5c70869 100644 --- a/clang/lib/AST/DeclCXX.cpp +++ b/clang/lib/AST/DeclCXX.cpp @@ -2005,7 +2005,9 @@ CXXMethodDecl *CXXMethodDecl::getDevirtualizedMethod(const Expr *Base, return nullptr; } -bool CXXMethodDecl::isUsualDeallocationFunction() const { +bool CXXMethodDecl::isUsualDeallocationFunction( + SmallVectorImpl &PreventedBy) const { + assert(PreventedBy.empty() && "PreventedBy is expected to be empty"); if (getOverloadedOperator() != OO_Delete && getOverloadedOperator() != OO_Array_Delete) return false; @@ -2063,14 +2065,16 @@ bool CXXMethodDecl::isUsualDeallocationFunction() const { // This function is a usual deallocation function if there are no // single-parameter deallocation functions of the same kind. DeclContext::lookup_result R = getDeclContext()->lookup(getDeclName()); - for (DeclContext::lookup_result::iterator I = R.begin(), E = R.end(); - I != E; ++I) { - if (const auto *FD = dyn_cast(*I)) - if (FD->getNumParams() == 1) - return false; + bool Result = true; + for (const auto *D : R) { + if (const auto *FD = dyn_cast(D)) { + if (FD->getNumParams() == 1) { + PreventedBy.push_back(FD); + Result = false; + } + } } - - return true; + return Result; } bool CXXMethodDecl::isCopyAssignmentOperator() const { diff --git a/clang/lib/Sema/SemaDeclCXX.cpp b/clang/lib/Sema/SemaDeclCXX.cpp index 95ad45d..eb3a2fc 100644 --- a/clang/lib/Sema/SemaDeclCXX.cpp +++ b/clang/lib/Sema/SemaDeclCXX.cpp @@ -13183,7 +13183,8 @@ CheckOperatorDeleteDeclaration(Sema &SemaRef, FunctionDecl *FnDecl) { // C++ P0722: // A destroying operator delete shall be a usual deallocation function. if (MD && !MD->getParent()->isDependentContext() && - MD->isDestroyingOperatorDelete() && !MD->isUsualDeallocationFunction()) { + MD->isDestroyingOperatorDelete() && + !SemaRef.isUsualDeallocationFunction(MD)) { SemaRef.Diag(MD->getLocation(), diag::err_destroying_operator_delete_not_usual); return true; diff --git a/clang/lib/Sema/SemaExprCXX.cpp b/clang/lib/Sema/SemaExprCXX.cpp index 1831430..dd5dfaa 100644 --- a/clang/lib/Sema/SemaExprCXX.cpp +++ b/clang/lib/Sema/SemaExprCXX.cpp @@ -1448,11 +1448,33 @@ Sema::BuildCXXTypeConstructExpr(TypeSourceInfo *TInfo, return Result; } +bool Sema::isUsualDeallocationFunction(const CXXMethodDecl *Method) { + // [CUDA] Ignore this function, if we can't call it. + const FunctionDecl *Caller = dyn_cast(CurContext); + if (getLangOpts().CUDA && + IdentifyCUDAPreference(Caller, Method) <= CFP_WrongSide) + return false; + + SmallVector PreventedBy; + bool Result = Method->isUsualDeallocationFunction(PreventedBy); + + if (Result || !getLangOpts().CUDA || PreventedBy.empty()) + return Result; + + // In case of CUDA, return true if none of the 1-argument deallocator + // functions are actually callable. + return llvm::none_of(PreventedBy, [&](const FunctionDecl *FD) { + assert(FD->getNumParams() == 1 && + "Only single-operand functions should be in PreventedBy"); + return IdentifyCUDAPreference(Caller, FD) >= CFP_HostDevice; + }); +} + /// Determine whether the given function is a non-placement /// deallocation function. static bool isNonPlacementDeallocationFunction(Sema &S, FunctionDecl *FD) { if (CXXMethodDecl *Method = dyn_cast(FD)) - return Method->isUsualDeallocationFunction(); + return S.isUsualDeallocationFunction(Method); if (FD->getOverloadedOperator() != OO_Delete && FD->getOverloadedOperator() != OO_Array_Delete) diff --git a/clang/test/CodeGenCUDA/usual-deallocators.cu b/clang/test/CodeGenCUDA/usual-deallocators.cu new file mode 100644 index 0000000..2d97c8c --- /dev/null +++ b/clang/test/CodeGenCUDA/usual-deallocators.cu @@ -0,0 +1,133 @@ +// RUN: %clang_cc1 %s --std=c++11 -triple nvptx-unknown-unknown -fcuda-is-device \ +// RUN: -emit-llvm -o - | FileCheck %s --check-prefixes=COMMON,DEVICE +// RUN: %clang_cc1 %s --std=c++11 -triple nvptx-unknown-unknown \ +// RUN: -emit-llvm -o - | FileCheck %s --check-prefixes=COMMON,HOST +// RUN: %clang_cc1 %s --std=c++17 -triple nvptx-unknown-unknown -fcuda-is-device \ +// RUN: -emit-llvm -o - | FileCheck %s --check-prefixes=COMMON,DEVICE +// RUN: %clang_cc1 %s --std=c++17 -triple nvptx-unknown-unknown \ +// RUN: -emit-llvm -o - | FileCheck %s --check-prefixes=COMMON,HOST + +#include "Inputs/cuda.h" +extern "C" __host__ void host_fn(); +extern "C" __device__ void dev_fn(); +extern "C" __host__ __device__ void hd_fn(); + +struct H1D1 { + __host__ void operator delete(void *) { host_fn(); }; + __device__ void operator delete(void *) { dev_fn(); }; +}; + +struct H1D2 { + __host__ void operator delete(void *) { host_fn(); }; + __device__ void operator delete(void *, __SIZE_TYPE__) { dev_fn(); }; +}; + +struct H2D1 { + __host__ void operator delete(void *, __SIZE_TYPE__) { host_fn(); }; + __device__ void operator delete(void *) { dev_fn(); }; +}; + +struct H2D2 { + __host__ void operator delete(void *, __SIZE_TYPE__) { host_fn(); }; + __device__ void operator delete(void *, __SIZE_TYPE__) { dev_fn(); }; +}; + +struct H1D1D2 { + __host__ void operator delete(void *) { host_fn(); }; + __device__ void operator delete(void *) { dev_fn(); }; + __device__ void operator delete(void *, __SIZE_TYPE__) { dev_fn(); }; +}; + +struct H1H2D1 { + __host__ void operator delete(void *) { host_fn(); }; + __host__ void operator delete(void *, __SIZE_TYPE__) { host_fn(); }; + __device__ void operator delete(void *) { dev_fn(); }; +}; + +struct H1H2D2 { + __host__ void operator delete(void *) { host_fn(); }; + __host__ void operator delete(void *, __SIZE_TYPE__) { host_fn(); }; + __device__ void operator delete(void *, __SIZE_TYPE__) { dev_fn(); }; +}; + +struct H1H2D1D2 { + __host__ void operator delete(void *) { host_fn(); }; + __host__ void operator delete(void *, __SIZE_TYPE__) { host_fn(); }; + __device__ void operator delete(void *) { dev_fn(); }; + __device__ void operator delete(void *, __SIZE_TYPE__) { dev_fn(); }; +}; + + +template +__host__ __device__ void test_hd(void *p) { + T *t = (T *)p; + delete t; +} + +// Make sure we call the right variant of usual deallocator. +__host__ __device__ void tests_hd(void *t) { + // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI4H1D1EvPv + // COMMON: call void @_ZN4H1D1dlEPv + test_hd(t); + // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI4H1D2EvPv + // DEVICE: call void @_ZN4H1D2dlEPvj(i8* {{.*}}, i32 1) + // HOST: call void @_ZN4H1D2dlEPv(i8* {{.*}}) + test_hd(t); + // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI4H2D1EvPv + // DEVICE: call void @_ZN4H2D1dlEPv(i8* {{.*}}) + // HOST: call void @_ZN4H2D1dlEPvj(i8* %3, i32 1) + test_hd(t); + // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI4H2D2EvPv + // COMMON: call void @_ZN4H2D2dlEPvj(i8* {{.*}}, i32 1) + test_hd(t); + // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI6H1D1D2EvPv + // COMMON: call void @_ZN6H1D1D2dlEPv(i8* %3) + test_hd(t); + // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI6H1H2D1EvPv + // COMMON: call void @_ZN6H1H2D1dlEPv(i8* {{.*}}) + test_hd(t); + // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI6H1H2D2EvPv + // DEVICE: call void @_ZN6H1H2D2dlEPvj(i8* {{.*}}, i32 1) + // HOST: call void @_ZN6H1H2D2dlEPv(i8* {{.*}}) + test_hd(t); + // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI8H1H2D1D2EvPv + // COMMON: call void @_ZN8H1H2D1D2dlEPv(i8* {{.*}}) + test_hd(t); +} + +// Make sure we've picked deallocator for the correct side of compilation. + +// COMMON-LABEL: define linkonce_odr void @_ZN4H1D1dlEPv(i8*) +// DEVICE: call void @dev_fn() +// HOST: call void @host_fn() + +// DEVICE-LABEL: define linkonce_odr void @_ZN4H1D2dlEPvj(i8*, i32) +// DEVICE: call void @dev_fn() +// HOST-LABEL: define linkonce_odr void @_ZN4H1D2dlEPv(i8*) +// HOST: call void @host_fn() + +// DEVICE-LABEL: define linkonce_odr void @_ZN4H2D1dlEPv(i8*) +// DEVICE: call void @dev_fn() +// HOST-LABEL: define linkonce_odr void @_ZN4H2D1dlEPvj(i8*, i32) +// HOST: call void @host_fn() + +// COMMON-LABEL: define linkonce_odr void @_ZN4H2D2dlEPvj(i8*, i32) +// DEVICE: call void @dev_fn() +// HOST: call void @host_fn() + +// COMMON-LABEL: define linkonce_odr void @_ZN6H1D1D2dlEPv(i8*) +// DEVICE: call void @dev_fn() +// HOST: call void @host_fn() + +// COMMON-LABEL: define linkonce_odr void @_ZN6H1H2D1dlEPv(i8*) +// DEVICE: call void @dev_fn() +// HOST: call void @host_fn() + +// DEVICE-LABEL: define linkonce_odr void @_ZN6H1H2D2dlEPvj(i8*, i32) +// DEVICE: call void @dev_fn() +// HOST-LABEL: define linkonce_odr void @_ZN6H1H2D2dlEPv(i8*) +// HOST: call void @host_fn() + +// COMMON-LABEL: define linkonce_odr void @_ZN8H1H2D1D2dlEPv(i8*) +// DEVICE: call void @dev_fn() +// HOST: call void @host_fn() diff --git a/clang/test/SemaCUDA/call-host-fn-from-device.cu b/clang/test/SemaCUDA/call-host-fn-from-device.cu index c727dc8..acdd291 100644 --- a/clang/test/SemaCUDA/call-host-fn-from-device.cu +++ b/clang/test/SemaCUDA/call-host-fn-from-device.cu @@ -41,12 +41,12 @@ struct T { operator Dummy() { return Dummy(); } // expected-note@-1 {{'operator Dummy' declared here}} - __host__ void operator delete(void*); - __device__ void operator delete(void*, size_t); + __host__ void operator delete(void *) { host_fn(); }; + __device__ void operator delete(void*, __SIZE_TYPE__); }; struct U { - __device__ void operator delete(void*, size_t) = delete; + __device__ void operator delete(void*, __SIZE_TYPE__) = delete; __host__ __device__ void operator delete(void*); }; diff --git a/clang/test/SemaCUDA/usual-deallocators.cu b/clang/test/SemaCUDA/usual-deallocators.cu new file mode 100644 index 0000000..a023864 --- /dev/null +++ b/clang/test/SemaCUDA/usual-deallocators.cu @@ -0,0 +1,95 @@ +// RUN: %clang_cc1 %s --std=c++11 -triple nvptx-unknown-unknown -fcuda-is-device \ +// RUN: -emit-llvm -o /dev/null -verify=device +// RUN: %clang_cc1 %s --std=c++11 -triple nvptx-unknown-unknown \ +// RUN: -emit-llvm -o /dev/null -verify=host +// RUN: %clang_cc1 %s --std=c++17 -triple nvptx-unknown-unknown -fcuda-is-device \ +// RUN: -emit-llvm -o /dev/null -verify=device +// RUN: %clang_cc1 %s --std=c++17 -triple nvptx-unknown-unknown \ +// RUN: -emit-llvm -o /dev/null -verify=host + +#include "Inputs/cuda.h" +extern __host__ void host_fn(); +extern __device__ void dev_fn(); +extern __host__ __device__ void hd_fn(); + +struct H1D1 { + __host__ void operator delete(void *) { host_fn(); }; + __device__ void operator delete(void *) { dev_fn(); }; +}; + +struct h1D1 { + __host__ void operator delete(void *) = delete; + // host-note@-1 {{'operator delete' has been explicitly marked deleted here}} + __device__ void operator delete(void *) { dev_fn(); }; +}; + +struct H1d1 { + __host__ void operator delete(void *) { host_fn(); }; + __device__ void operator delete(void *) = delete; + // device-note@-1 {{'operator delete' has been explicitly marked deleted here}} +}; + +struct H1D2 { + __host__ void operator delete(void *) { host_fn(); }; + __device__ void operator delete(void *, __SIZE_TYPE__) { dev_fn(); }; +}; + +struct H2D1 { + __host__ void operator delete(void *, __SIZE_TYPE__) { host_fn(); }; + __device__ void operator delete(void *) { dev_fn(); }; +}; + +struct H2D2 { + __host__ void operator delete(void *, __SIZE_TYPE__) { host_fn(); }; + __device__ void operator delete(void *, __SIZE_TYPE__) { dev_fn(); }; +}; + +struct H1D1D2 { + __host__ void operator delete(void *) { host_fn(); }; + __device__ void operator delete(void *) { dev_fn(); }; + __device__ void operator delete(void *, __SIZE_TYPE__) { dev_fn(); }; +}; + +struct H1H2D1 { + __host__ void operator delete(void *) { host_fn(); }; + __host__ void operator delete(void *, __SIZE_TYPE__) { host_fn(); }; + __device__ void operator delete(void *) { dev_fn(); }; +}; + +struct H1H2D2 { + __host__ void operator delete(void *) { host_fn(); }; + __host__ void operator delete(void *, __SIZE_TYPE__) { host_fn(); }; + __device__ void operator delete(void *, __SIZE_TYPE__) { dev_fn(); }; +}; + +struct H1H2D1D2 { + __host__ void operator delete(void *) { host_fn(); }; + __host__ void operator delete(void *, __SIZE_TYPE__) { host_fn(); }; + __device__ void operator delete(void *) { dev_fn(); }; + __device__ void operator delete(void *, __SIZE_TYPE__) { dev_fn(); }; +}; + + +template +__host__ __device__ void test_hd(void *p) { + T *t = (T *)p; + delete t; + // host-error@-1 {{attempt to use a deleted function}} + // device-error@-2 {{attempt to use a deleted function}} +} + +__host__ __device__ void tests_hd(void *t) { + test_hd(t); + test_hd(t); + // host-note@-1 {{in instantiation of function template specialization 'test_hd' requested here}} + test_hd(t); + // device-note@-1 {{in instantiation of function template specialization 'test_hd' requested here}} + test_hd(t); + test_hd(t); + test_hd(t); + test_hd(t); + test_hd(t); + test_hd(t); + test_hd(t); + test_hd(t); +} -- 2.7.4