70 lines · plain
1// expected-no-diagnostics2 3// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only -verify %s4// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fsyntax-only -fcuda-is-device -verify %s5 6#include "Inputs/cuda.h"7 8struct S {9 __host__ static void operator delete(void*, size_t) {}10 __device__ static void operator delete(void*, size_t) {}11};12 13__host__ __device__ void test(S* s) {14 // This shouldn't be ambiguous -- we call the host overload in host mode and15 // the device overload in device mode.16 delete s;17}18 19// Code should work with no explicit declarations/definitions of20// allocator functions.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 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}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 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 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}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.54__host__ void operator delete(void *ptr) {}55__device__ void operator delete(void *ptr) {}56 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 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}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