This revision was automatically updated to reflect the committed changes.
Closed by commit rG247cc265e74e: [CUDA][HIP] Fix overloading resolution of
delete operator (authored by yaxunl).
Herald added a project: clang.
Changed prior to commit:
https://reviews.llvm.org/D156795?vs=546045&id=548198#toc
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D156795/new/
https://reviews.llvm.org/D156795
Files:
clang/lib/Sema/SemaExprCXX.cpp
clang/test/CodeGenCUDA/member-init.cu
clang/test/SemaCUDA/member-init.cu
Index: clang/test/SemaCUDA/member-init.cu
===================================================================
--- /dev/null
+++ 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;
+}
Index: clang/test/CodeGenCUDA/member-init.cu
===================================================================
--- /dev/null
+++ 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;
+}
Index: clang/lib/Sema/SemaExprCXX.cpp
===================================================================
--- clang/lib/Sema/SemaExprCXX.cpp
+++ clang/lib/Sema/SemaExprCXX.cpp
@@ -1714,8 +1714,8 @@
// 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; }
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
