1 // expected-no-diagnostics
2
3 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only -verify %s
4 // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fsyntax-only -fcuda-is-device -verify %s
5
6 #include "Inputs/cuda.h"
7
8 struct S {
operator deleteS9 __host__ static void operator delete(void*, size_t) {}
operator deleteS10 __device__ static void operator delete(void*, size_t) {}
11 };
12
test(S * s)13 __host__ __device__ void test(S* s) {
14 // This shouldn't be ambiguous -- we call the host overload in host mode and
15 // the device overload in device mode.
16 delete s;
17 }
18
19 // Code should work with no explicit declarations/definitions of
20 // allocator functions.
test_default_global_delete_hd(int * ptr)21 __host__ __device__ void test_default_global_delete_hd(int *ptr) {
22 // Again, there should be no ambiguity between which operator delete we call.
23 ::delete ptr;
24 }
25
test_default_global_delete(int * ptr)26 __device__ void test_default_global_delete(int *ptr) {
27 // Again, there should be no ambiguity between which operator delete we call.
28 ::delete ptr;
29 }
test_default_global_delete(int * ptr)30 __host__ void test_default_global_delete(int *ptr) {
31 // Again, there should be no ambiguity between which operator delete we call.
32 ::delete ptr;
33 }
34
35 // It should work with only some of allocators (re-)declared.
36 __device__ void operator delete(void *ptr);
37
test_partial_global_delete_hd(int * ptr)38 __host__ __device__ void test_partial_global_delete_hd(int *ptr) {
39 // Again, there should be no ambiguity between which operator delete we call.
40 ::delete ptr;
41 }
42
test_partial_global_delete(int * ptr)43 __device__ void test_partial_global_delete(int *ptr) {
44 // Again, there should be no ambiguity between which operator delete we call.
45 ::delete ptr;
46 }
test_partial_global_delete(int * ptr)47 __host__ void test_partial_global_delete(int *ptr) {
48 // Again, there should be no ambiguity between which operator delete we call.
49 ::delete ptr;
50 }
51
52
53 // We should be able to define both host and device variants.
operator delete(void * ptr)54 __host__ void operator delete(void *ptr) {}
operator delete(void * ptr)55 __device__ void operator delete(void *ptr) {}
56
test_overloaded_global_delete_hd(int * ptr)57 __host__ __device__ void test_overloaded_global_delete_hd(int *ptr) {
58 // Again, there should be no ambiguity between which operator delete we call.
59 ::delete ptr;
60 }
61
test_overloaded_global_delete(int * ptr)62 __device__ void test_overloaded_global_delete(int *ptr) {
63 // Again, there should be no ambiguity between which operator delete we call.
64 ::delete ptr;
65 }
test_overloaded_global_delete(int * ptr)66 __host__ void test_overloaded_global_delete(int *ptr) {
67 // Again, there should be no ambiguity between which operator delete we call.
68 ::delete ptr;
69 }
70