~shimon/dismaltc

ref: 461e0617cb03f1d5027af4fbb3da0722791022e9 dismaltc/clang/disable-test/CodeGenCUDA/usual-deallocators.cu -rw-r--r-- 4.9 KiB
461e0617Itoh Shimon re-enable tests. 1 year, 6 months ago
                                                                                
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
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 <typename T>
__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<H1D1>(t);
  // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI4H1D2EvPv
  // DEVICE: call void @_ZN4H1D2dlEPvj(i8* {{.*}}, i32 1)
  // HOST:   call void @_ZN4H1D2dlEPv(i8* {{.*}})
  test_hd<H1D2>(t);
  // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI4H2D1EvPv
  // DEVICE: call void @_ZN4H2D1dlEPv(i8* {{.*}})
  // HOST:   call void @_ZN4H2D1dlEPvj(i8* %3, i32 1)
  test_hd<H2D1>(t);
  // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI4H2D2EvPv
  // COMMON: call void @_ZN4H2D2dlEPvj(i8* {{.*}}, i32 1)
  test_hd<H2D2>(t);
  // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI6H1D1D2EvPv
  // COMMON: call void @_ZN6H1D1D2dlEPv(i8* %3)
  test_hd<H1D1D2>(t);
  // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI6H1H2D1EvPv
  // COMMON: call void @_ZN6H1H2D1dlEPv(i8* {{.*}})
  test_hd<H1H2D1>(t);
  // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI6H1H2D2EvPv
  // DEVICE: call void @_ZN6H1H2D2dlEPvj(i8* {{.*}}, i32 1)
  // HOST:   call void @_ZN6H1H2D2dlEPv(i8* {{.*}})
  test_hd<H1H2D2>(t);
  // COMMON-LABEL: define linkonce_odr void @_Z7test_hdI8H1H2D1D2EvPv
  // COMMON: call void @_ZN8H1H2D1D2dlEPv(i8* {{.*}})
  test_hd<H1H2D1D2>(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()