1 // RUN: %clang_cc1 -triple x86_64-pc-linux-gnu -emit-llvm -fexceptions \
2 // RUN: -o - -x hip %s | FileCheck %s
3
4 #include "Inputs/cuda.h"
5
6 int* hvar;
7 __device__ int* dvar;
8
9 // CHECK-LABEL: define {{.*}}@_Znwm
10 // CHECK: load ptr, ptr @hvar
operator new(unsigned long size)11 void* operator new(unsigned long size) {
12 return hvar;
13 }
14 // CHECK-LABEL: define {{.*}}@_ZdlPv
15 // CHECK: store ptr inttoptr (i64 1 to ptr), ptr @hvar
operator delete(void * p)16 void operator delete(void *p) {
17 hvar = (int*)1;
18 }
19
operator new(unsigned long size)20 __device__ void* operator new(unsigned long size) {
21 return dvar;
22 }
23
operator delete(void * p)24 __device__ void operator delete(void *p) {
25 dvar = (int*)11;
26 }
27
28 class A {
29 int x;
30 public:
A()31 A(){
32 x = 123;
33 }
34 };
35
36 template<class T>
37 class shared_ptr {
38 int id;
39 T *ptr;
40 public:
shared_ptr(T * p)41 shared_ptr(T *p) {
42 id = 2;
43 ptr = p;
44 }
45 };
46
47 // The constructor of B calls the delete operator to clean up
48 // the memory allocated by the new operator when exceptions happen.
49 // Make sure the host delete operator is used on host side.
50 //
51 // No need to do similar checks on the device side since it does
52 // not support exception.
53
54 // CHECK-LABEL: define {{.*}}@main
55 // CHECK: call void @_ZN1BC1Ev
56
57 // CHECK-LABEL: define {{.*}}@_ZN1BC1Ev
58 // CHECK: call void @_ZN1BC2Ev
59
60 // CHECK-LABEL: define {{.*}}@_ZN1BC2Ev
61 // CHECK: call {{.*}}@_Znwm
62 // CHECK: invoke void @_ZN1AC1Ev
63 // CHECK: call void @_ZN10shared_ptrI1AEC1EPS0_
64 // CHECK: cleanup
65 // CHECK: call void @_ZdlPv
66
67 struct B{
68 shared_ptr<A> pa{new A};
69 };
70
main()71 int main() {
72 B b;
73 }
74