| // 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; | |
| } |