diff options
Diffstat (limited to 'clang/test/SemaCUDA/function-overload.cu')
| -rw-r--r-- | clang/test/SemaCUDA/function-overload.cu | 317 |
1 files changed, 317 insertions, 0 deletions
diff --git a/clang/test/SemaCUDA/function-overload.cu b/clang/test/SemaCUDA/function-overload.cu new file mode 100644 index 00000000000..bd3fb508bfa --- /dev/null +++ b/clang/test/SemaCUDA/function-overload.cu @@ -0,0 +1,317 @@ +// REQUIRES: x86-registered-target +// REQUIRES: nvptx-registered-target + +// Make sure we handle target overloads correctly. +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu \ +// RUN: -fsyntax-only -fcuda-target-overloads -verify %s +// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda \ +// RUN: -fsyntax-only -fcuda-target-overloads -fcuda-is-device -verify %s + +// Check target overloads handling with disabled call target checks. +// RUN: %clang_cc1 -DNOCHECKS -triple x86_64-unknown-linux-gnu -fsyntax-only \ +// RUN: -fcuda-disable-target-call-checks -fcuda-target-overloads -verify %s +// RUN: %clang_cc1 -DNOCHECKS -triple nvptx64-nvidia-cuda -fsyntax-only \ +// RUN: -fcuda-disable-target-call-checks -fcuda-target-overloads \ +// RUN: -fcuda-is-device -verify %s + +#include "Inputs/cuda.h" + +typedef int (*fp_t)(void); +typedef void (*gp_t)(void); + +// Host and unattributed functions can't be overloaded +__host__ int hh(void) { return 1; } // expected-note {{previous definition is here}} +int hh(void) { return 1; } // expected-error {{redefinition of 'hh'}} + +// H/D overloading is OK +__host__ int dh(void) { return 2; } +__device__ int dh(void) { return 2; } + +// H/HD and D/HD are not allowed +__host__ __device__ int hdh(void) { return 5; } // expected-note {{previous definition is here}} +__host__ int hdh(void) { return 4; } // expected-error {{redefinition of 'hdh'}} + +__host__ int hhd(void) { return 4; } // expected-note {{previous definition is here}} +__host__ __device__ int hhd(void) { return 5; } // expected-error {{redefinition of 'hhd'}} +// expected-warning@-1 {{attribute declaration must precede definition}} +// expected-note@-3 {{previous definition is here}} + +__host__ __device__ int hdd(void) { return 7; } // expected-note {{previous definition is here}} +__device__ int hdd(void) { return 6; } // expected-error {{redefinition of 'hdd'}} + +__device__ int dhd(void) { return 6; } // expected-note {{previous definition is here}} +__host__ __device__ int dhd(void) { return 7; } // expected-error {{redefinition of 'dhd'}} +// expected-warning@-1 {{attribute declaration must precede definition}} +// expected-note@-3 {{previous definition is here}} + +// Same tests for extern "C" functions +extern "C" __host__ int chh(void) {return 11;} // expected-note {{previous definition is here}} +extern "C" int chh(void) {return 11;} // expected-error {{redefinition of 'chh'}} + +// H/D overloading is OK +extern "C" __device__ int cdh(void) {return 10;} +extern "C" __host__ int cdh(void) {return 11;} + +// H/HD and D/HD overloading is not allowed. +extern "C" __host__ __device__ int chhd1(void) {return 12;} // expected-note {{previous definition is here}} +extern "C" __host__ int chhd1(void) {return 13;} // expected-error {{redefinition of 'chhd1'}} + +extern "C" __host__ int chhd2(void) {return 13;} // expected-note {{previous definition is here}} +extern "C" __host__ __device__ int chhd2(void) {return 12;} // expected-error {{redefinition of 'chhd2'}} +// expected-warning@-1 {{attribute declaration must precede definition}} +// expected-note@-3 {{previous definition is here}} + +// Helper functions to verify calling restrictions. +__device__ int d(void) { return 8; } +__host__ int h(void) { return 9; } +__global__ void g(void) {} +extern "C" __device__ int cd(void) {return 10;} +extern "C" __host__ int ch(void) {return 11;} + +__host__ void hostf(void) { + fp_t dp = d; + fp_t cdp = cd; +#if !defined(NOCHECKS) + // expected-error@-3 {{reference to __device__ function 'd' in __host__ function}} + // expected-note@65 {{'d' declared here}} + // expected-error@-4 {{reference to __device__ function 'cd' in __host__ function}} + // expected-note@68 {{'cd' declared here}} +#endif + fp_t hp = h; + fp_t chp = ch; + fp_t dhp = dh; + fp_t cdhp = cdh; + gp_t gp = g; + + d(); + cd(); +#if !defined(NOCHECKS) + // expected-error@-3 {{no matching function for call to 'd'}} + // expected-note@65 {{candidate function not viable: call to __device__ function from __host__ function}} + // expected-error@-4 {{no matching function for call to 'cd'}} + // expected-note@68 {{candidate function not viable: call to __device__ function from __host__ function}} +#endif + h(); + ch(); + dh(); + cdh(); + g(); // expected-error {{call to global function g not configured}} + g<<<0,0>>>(); +} + + +__device__ void devicef(void) { + fp_t dp = d; + fp_t cdp = cd; + fp_t hp = h; + fp_t chp = ch; +#if !defined(NOCHECKS) + // expected-error@-3 {{reference to __host__ function 'h' in __device__ function}} + // expected-note@66 {{'h' declared here}} + // expected-error@-4 {{reference to __host__ function 'ch' in __device__ function}} + // expected-note@69 {{'ch' declared here}} +#endif + fp_t dhp = dh; + fp_t cdhp = cdh; + gp_t gp = g; // expected-error {{reference to __global__ function 'g' in __device__ function}} + // expected-note@67 {{'g' declared here}} + + d(); + cd(); + h(); + ch(); +#if !defined(NOCHECKS) + // expected-error@-3 {{no matching function for call to 'h'}} + // expected-note@66 {{candidate function not viable: call to __host__ function from __device__ function}} + // expected-error@-4 {{no matching function for call to 'ch'}} + // expected-note@69 {{candidate function not viable: call to __host__ function from __device__ function}} +#endif + dh(); + cdh(); + g(); // expected-error {{no matching function for call to 'g'}} + // expected-note@67 {{candidate function not viable: call to __global__ function from __device__ function}} + g<<<0,0>>>(); // expected-error {{reference to __global__ function 'g' in __device__ function}} + // expected-note@67 {{'g' declared here}} +} + +__global__ void globalf(void) { + fp_t dp = d; + fp_t cdp = cd; + fp_t hp = h; + fp_t chp = ch; +#if !defined(NOCHECKS) + // expected-error@-3 {{reference to __host__ function 'h' in __global__ function}} + // expected-note@66 {{'h' declared here}} + // expected-error@-4 {{reference to __host__ function 'ch' in __global__ function}} + // expected-note@69 {{'ch' declared here}} +#endif + fp_t dhp = dh; + fp_t cdhp = cdh; + gp_t gp = g; // expected-error {{reference to __global__ function 'g' in __global__ function}} + // expected-note@67 {{'g' declared here}} + + d(); + cd(); + h(); + ch(); +#if !defined(NOCHECKS) + // expected-error@-3 {{no matching function for call to 'h'}} + // expected-note@66 {{candidate function not viable: call to __host__ function from __global__ function}} + // expected-error@-4 {{no matching function for call to 'ch'}} + // expected-note@69 {{candidate function not viable: call to __host__ function from __global__ function}} +#endif + dh(); + cdh(); + g(); // expected-error {{no matching function for call to 'g'}} + // expected-note@67 {{candidate function not viable: call to __global__ function from __global__ function}} + g<<<0,0>>>(); // expected-error {{reference to __global__ function 'g' in __global__ function}} + // expected-note@67 {{'g' declared here}} +} + +__host__ __device__ void hostdevicef(void) { + fp_t dp = d; + fp_t cdp = cd; + fp_t hp = h; + fp_t chp = ch; +#if !defined(NOCHECKS) +#if !defined(__CUDA_ARCH__) + // expected-error@-6 {{reference to __device__ function 'd' in __host__ __device__ function}} + // expected-note@65 {{'d' declared here}} + // expected-error@-7 {{reference to __device__ function 'cd' in __host__ __device__ function}} + // expected-note@68 {{'cd' declared here}} +#else + // expected-error@-9 {{reference to __host__ function 'h' in __host__ __device__ function}} + // expected-note@66 {{'h' declared here}} + // expected-error@-10 {{reference to __host__ function 'ch' in __host__ __device__ function}} + // expected-note@69 {{'ch' declared here}} +#endif +#endif + fp_t dhp = dh; + fp_t cdhp = cdh; + gp_t gp = g; +#if defined(__CUDA_ARCH__) + // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}} + // expected-note@67 {{'g' declared here}} +#endif + + d(); + cd(); + h(); + ch(); +#if !defined(NOCHECKS) +#if !defined(__CUDA_ARCH__) + // expected-error@-6 {{no matching function for call to 'd'}} + // expected-note@65 {{candidate function not viable: call to __device__ function from __host__ __device__ function}} + // expected-error@-7 {{no matching function for call to 'cd'}} + // expected-note@68 {{candidate function not viable: call to __device__ function from __host__ __device__ function}} +#else + // expected-error@-9 {{no matching function for call to 'h'}} + // expected-note@66 {{candidate function not viable: call to __host__ function from __host__ __device__ function}} + // expected-error@-10 {{no matching function for call to 'ch'}} + // expected-note@69 {{candidate function not viable: call to __host__ function from __host__ __device__ function}} +#endif +#endif + + dh(); + cdh(); + g(); + g<<<0,0>>>(); +#if !defined(__CUDA_ARCH__) + // expected-error@-3 {{call to global function g not configured}} +#else + // expected-error@-5 {{no matching function for call to 'g'}} + // expected-note@67 {{candidate function not viable: call to __global__ function from __host__ __device__ function}} + // expected-error@-6 {{reference to __global__ function 'g' in __host__ __device__ function}} + // expected-note@67 {{'g' declared here}} +#endif // __CUDA_ARCH__ +} + +// Test for address of overloaded function resolution in the global context. +fp_t hp = h; +fp_t chp = ch; +fp_t dhp = dh; +fp_t cdhp = cdh; +gp_t gp = g; + + +// Test overloading of destructors +// Can't mix H and unattributed destructors +struct d_h { + ~d_h() {} // expected-note {{previous declaration is here}} + __host__ ~d_h() {} // expected-error {{destructor cannot be redeclared}} +}; + +// H/D overloading is OK +struct d_dh { + __device__ ~d_dh() {} + __host__ ~d_dh() {} +}; + +// HD is OK +struct d_hd { + __host__ __device__ ~d_hd() {} +}; + +// Mixing H/D and HD is not allowed. +struct d_dhhd { + __device__ ~d_dhhd() {} + __host__ ~d_dhhd() {} // expected-note {{previous declaration is here}} + __host__ __device__ ~d_dhhd() {} // expected-error {{destructor cannot be redeclared}} +}; + +struct d_hhd { + __host__ ~d_hhd() {} // expected-note {{previous declaration is here}} + __host__ __device__ ~d_hhd() {} // expected-error {{destructor cannot be redeclared}} +}; + +struct d_hdh { + __host__ __device__ ~d_hdh() {} // expected-note {{previous declaration is here}} + __host__ ~d_hdh() {} // expected-error {{destructor cannot be redeclared}} +}; + +struct d_dhd { + __device__ ~d_dhd() {} // expected-note {{previous declaration is here}} + __host__ __device__ ~d_dhd() {} // expected-error {{destructor cannot be redeclared}} +}; + +struct d_hdd { + __host__ __device__ ~d_hdd() {} // expected-note {{previous declaration is here}} + __device__ ~d_hdd() {} // expected-error {{destructor cannot be redeclared}} +}; + +// Test overloading of member functions +struct m_h { + void operator delete(void *ptr); // expected-note {{previous declaration is here}} + __host__ void operator delete(void *ptr); // expected-error {{class member cannot be redeclared}} +}; + +// D/H overloading is OK +struct m_dh { + __device__ void operator delete(void *ptr); + __host__ void operator delete(void *ptr); +}; + +// HD by itself is OK +struct m_hd { + __device__ __host__ void operator delete(void *ptr); +}; + +struct m_hhd { + __host__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}} + __host__ __device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}} +}; + +struct m_hdh { + __host__ __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}} + __host__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}} +}; + +struct m_dhd { + __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}} + __host__ __device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}} +}; + +struct m_hdd { + __host__ __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}} + __device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}} +}; |
