Author: Yaxun (Sam) Liu Date: 2023-08-08T09:50:24-04:00 New Revision: 247cc265e74e25164ee7ce85e6c9c53b3d177740
URL: https://github.com/llvm/llvm-project/commit/247cc265e74e25164ee7ce85e6c9c53b3d177740 DIFF: https://github.com/llvm/llvm-project/commit/247cc265e74e25164ee7ce85e6c9c53b3d177740.diff LOG: [CUDA][HIP] Fix overloading resolution of delete operator Currently clang does not consider host/device preference when resolving delete operator in the file scope, which causes device operator delete selected for class member initialization. Reviewed by: Artem Belevich Differential Revision: https://reviews.llvm.org/D156795 Added: clang/test/CodeGenCUDA/member-init.cu clang/test/SemaCUDA/member-init.cu Modified: clang/lib/Sema/SemaExprCXX.cpp Removed: ################################################################################ diff --git a/clang/lib/Sema/SemaExprCXX.cpp b/clang/lib/Sema/SemaExprCXX.cpp index 423d5372a6f65a..82afb084efd0b1 100644 --- a/clang/lib/Sema/SemaExprCXX.cpp +++ b/clang/lib/Sema/SemaExprCXX.cpp @@ -1714,8 +1714,8 @@ namespace { // In CUDA, determine how much we'd like / dislike to call this. if (S.getLangOpts().CUDA) - if (auto *Caller = S.getCurFunctionDecl(/*AllowLambda=*/true)) - CUDAPref = S.IdentifyCUDAPreference(Caller, FD); + CUDAPref = S.IdentifyCUDAPreference( + S.getCurFunctionDecl(/*AllowLambda=*/true), FD); } explicit operator bool() const { return FD; } diff --git a/clang/test/CodeGenCUDA/member-init.cu b/clang/test/CodeGenCUDA/member-init.cu new file mode 100644 index 00000000000000..8d1db494a40e4a --- /dev/null +++ b/clang/test/CodeGenCUDA/member-init.cu @@ -0,0 +1,73 @@ +// RUN: %clang_cc1 -triple x86_64-pc-linux-gnu -emit-llvm -fexceptions \ +// RUN: -o - -x hip %s | FileCheck %s + +#include "Inputs/cuda.h" + +int* hvar; +__device__ int* dvar; + +// CHECK-LABEL: define {{.*}}@_Znwm +// CHECK: load ptr, ptr @hvar +void* operator new(unsigned long size) { + return hvar; +} +// CHECK-LABEL: define {{.*}}@_ZdlPv +// CHECK: store ptr inttoptr (i64 1 to ptr), ptr @hvar +void operator delete(void *p) { + hvar = (int*)1; +} + +__device__ void* operator new(unsigned long size) { + return dvar; +} + +__device__ void operator delete(void *p) { + dvar = (int*)11; +} + +class A { + int x; +public: + A(){ + x = 123; + } +}; + +template<class T> +class shared_ptr { + int id; + T *ptr; +public: + shared_ptr(T *p) { + id = 2; + ptr = p; + } +}; + +// The constructor of B calls the delete operator to clean up +// the memory allocated by the new operator when exceptions happen. +// Make sure the host delete operator is used on host side. +// +// No need to do similar checks on the device side since it does +// not support exception. + +// CHECK-LABEL: define {{.*}}@main +// CHECK: call void @_ZN1BC1Ev + +// CHECK-LABEL: define {{.*}}@_ZN1BC1Ev +// CHECK: call void @_ZN1BC2Ev + +// CHECK-LABEL: define {{.*}}@_ZN1BC2Ev +// CHECK: call {{.*}}@_Znwm +// CHECK: invoke void @_ZN1AC1Ev +// CHECK: call void @_ZN10shared_ptrI1AEC1EPS0_ +// CHECK: cleanup +// CHECK: call void @_ZdlPv + +struct B{ + shared_ptr<A> pa{new A}; +}; + +int main() { + B b; +} diff --git a/clang/test/SemaCUDA/member-init.cu b/clang/test/SemaCUDA/member-init.cu new file mode 100644 index 00000000000000..a796a50e9c8768 --- /dev/null +++ b/clang/test/SemaCUDA/member-init.cu @@ -0,0 +1,38 @@ +// RUN: %clang_cc1 -fsyntax-only -verify -fexceptions %s +// expected-no-diagnostics + +#include "Inputs/cuda.h" + +__device__ void operator delete(void *p) {} + +class A { + int x; +public: + A() { + x = 123; + } +}; + +template<class T> +class shared_ptr { + T *ptr; +public: + shared_ptr(T *p) { + ptr = p; + } +}; + +// The constructor of B calls the delete operator to clean up +// the memory allocated by the new operator when exceptions happen. +// Make sure that there are no diagnostics due to the device delete +// operator is used. +// +// No need to do similar checks on the device side since it does +// not support exception. +struct B{ + shared_ptr<A> pa{new A}; +}; + +int main() { + B b; +} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits