1 | // RUN: %clang_cc1 %s --std=c++11 -triple nvptx-unknown-unknown -fcuda-is-device \ |
2 | // RUN: -emit-llvm -o /dev/null -verify=device |
3 | // RUN: %clang_cc1 %s --std=c++11 -triple nvptx-unknown-unknown \ |
4 | // RUN: -emit-llvm -o /dev/null -verify=host |
5 | // RUN: %clang_cc1 %s --std=c++17 -triple nvptx-unknown-unknown -fcuda-is-device \ |
6 | // RUN: -emit-llvm -o /dev/null -verify=device |
7 | // RUN: %clang_cc1 %s --std=c++17 -triple nvptx-unknown-unknown \ |
8 | // RUN: -emit-llvm -o /dev/null -verify=host |
9 | |
10 | #include "Inputs/cuda.h" |
11 | extern __host__ void host_fn(); |
12 | extern __device__ void dev_fn(); |
13 | extern __host__ __device__ void hd_fn(); |
14 | |
15 | struct H1D1 { |
16 | __host__ void operator delete(void *) { host_fn(); }; |
17 | __device__ void operator delete(void *) { dev_fn(); }; |
18 | }; |
19 | |
20 | struct h1D1 { |
21 | __host__ void operator delete(void *) = delete; |
22 | // host-note@-1 {{'operator delete' has been explicitly marked deleted here}} |
23 | __device__ void operator delete(void *) { dev_fn(); }; |
24 | }; |
25 | |
26 | struct H1d1 { |
27 | __host__ void operator delete(void *) { host_fn(); }; |
28 | __device__ void operator delete(void *) = delete; |
29 | // device-note@-1 {{'operator delete' has been explicitly marked deleted here}} |
30 | }; |
31 | |
32 | struct H1D2 { |
33 | __host__ void operator delete(void *) { host_fn(); }; |
34 | __device__ void operator delete(void *, __SIZE_TYPE__) { dev_fn(); }; |
35 | }; |
36 | |
37 | struct H2D1 { |
38 | __host__ void operator delete(void *, __SIZE_TYPE__) { host_fn(); }; |
39 | __device__ void operator delete(void *) { dev_fn(); }; |
40 | }; |
41 | |
42 | struct H2D2 { |
43 | __host__ void operator delete(void *, __SIZE_TYPE__) { host_fn(); }; |
44 | __device__ void operator delete(void *, __SIZE_TYPE__) { dev_fn(); }; |
45 | }; |
46 | |
47 | struct H1D1D2 { |
48 | __host__ void operator delete(void *) { host_fn(); }; |
49 | __device__ void operator delete(void *) { dev_fn(); }; |
50 | __device__ void operator delete(void *, __SIZE_TYPE__) { dev_fn(); }; |
51 | }; |
52 | |
53 | struct H1H2D1 { |
54 | __host__ void operator delete(void *) { host_fn(); }; |
55 | __host__ void operator delete(void *, __SIZE_TYPE__) { host_fn(); }; |
56 | __device__ void operator delete(void *) { dev_fn(); }; |
57 | }; |
58 | |
59 | struct H1H2D2 { |
60 | __host__ void operator delete(void *) { host_fn(); }; |
61 | __host__ void operator delete(void *, __SIZE_TYPE__) { host_fn(); }; |
62 | __device__ void operator delete(void *, __SIZE_TYPE__) { dev_fn(); }; |
63 | }; |
64 | |
65 | struct H1H2D1D2 { |
66 | __host__ void operator delete(void *) { host_fn(); }; |
67 | __host__ void operator delete(void *, __SIZE_TYPE__) { host_fn(); }; |
68 | __device__ void operator delete(void *) { dev_fn(); }; |
69 | __device__ void operator delete(void *, __SIZE_TYPE__) { dev_fn(); }; |
70 | }; |
71 | |
72 | |
73 | template <typename T> |
74 | __host__ __device__ void test_hd(void *p) { |
75 | T *t = (T *)p; |
76 | delete t; |
77 | // host-error@-1 {{attempt to use a deleted function}} |
78 | // device-error@-2 {{attempt to use a deleted function}} |
79 | } |
80 | |
81 | __host__ __device__ void tests_hd(void *t) { |
82 | test_hd<H1D1>(t); |
83 | test_hd<h1D1>(t); |
84 | // host-note@-1 {{in instantiation of function template specialization 'test_hd<h1D1>' requested here}} |
85 | test_hd<H1d1>(t); |
86 | // device-note@-1 {{in instantiation of function template specialization 'test_hd<H1d1>' requested here}} |
87 | test_hd<H1D2>(t); |
88 | test_hd<H2D1>(t); |
89 | test_hd<H2D2>(t); |
90 | test_hd<H1D1D2>(t); |
91 | test_hd<H1H2D1>(t); |
92 | test_hd<H1H2D1>(t); |
93 | test_hd<H1H2D2>(t); |
94 | test_hd<H1H2D1D2>(t); |
95 | } |
96 | |