| // REQUIRES: x86-registered-target |
| // REQUIRES: nvptx-registered-target |
| |
| // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fsyntax-only -verify %s |
| // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fsyntax-only -fcuda-is-device -verify %s |
| |
| #include "Inputs/cuda.h" |
| |
| // Opaque return types used to check that we pick the right overloads. |
| struct HostReturnTy {}; |
| struct HostReturnTy2 {}; |
| struct DeviceReturnTy {}; |
| struct DeviceReturnTy2 {}; |
| struct HostDeviceReturnTy {}; |
| struct TemplateReturnTy {}; |
| |
| typedef HostReturnTy (*HostFnPtr)(); |
| typedef DeviceReturnTy (*DeviceFnPtr)(); |
| typedef HostDeviceReturnTy (*HostDeviceFnPtr)(); |
| typedef void (*GlobalFnPtr)(); // __global__ functions must return void. |
| |
| // CurrentReturnTy is {HostReturnTy,DeviceReturnTy} during {host,device} |
| // compilation. |
| #ifdef __CUDA_ARCH__ |
| typedef DeviceReturnTy CurrentReturnTy; |
| #else |
| typedef HostReturnTy CurrentReturnTy; |
| #endif |
| |
| // CurrentFnPtr is a function pointer to a {host,device} function during |
| // {host,device} compilation. |
| typedef CurrentReturnTy (*CurrentFnPtr)(); |
| |
| // Host and unattributed functions can't be overloaded. |
| __host__ void hh() {} // expected-note {{previous definition is here}} |
| void hh() {} // expected-error {{redefinition of 'hh'}} |
| |
| // H/D overloading is OK. |
| __host__ HostReturnTy dh() { return HostReturnTy(); } |
| __device__ DeviceReturnTy dh() { return DeviceReturnTy(); } |
| |
| // H/HD and D/HD are not allowed. |
| __host__ __device__ int hdh() { return 0; } // expected-note {{previous declaration is here}} |
| __host__ int hdh() { return 0; } |
| // expected-error@-1 {{__host__ function 'hdh' cannot overload __host__ __device__ function 'hdh'}} |
| |
| __host__ int hhd() { return 0; } // expected-note {{previous declaration is here}} |
| __host__ __device__ int hhd() { return 0; } |
| // expected-error@-1 {{__host__ __device__ function 'hhd' cannot overload __host__ function 'hhd'}} |
| |
| __host__ __device__ int hdd() { return 0; } // expected-note {{previous declaration is here}} |
| __device__ int hdd() { return 0; } |
| // expected-error@-1 {{__device__ function 'hdd' cannot overload __host__ __device__ function 'hdd'}} |
| |
| __device__ int dhd() { return 0; } // expected-note {{previous declaration is here}} |
| __host__ __device__ int dhd() { return 0; } |
| // expected-error@-1 {{__host__ __device__ function 'dhd' cannot overload __device__ function 'dhd'}} |
| |
| // Same tests for extern "C" functions. |
| extern "C" __host__ int chh() { return 0; } // expected-note {{previous definition is here}} |
| extern "C" int chh() { return 0; } // expected-error {{redefinition of 'chh'}} |
| |
| // H/D overloading is OK. |
| extern "C" __device__ DeviceReturnTy cdh() { return DeviceReturnTy(); } |
| extern "C" __host__ HostReturnTy cdh() { return HostReturnTy(); } |
| |
| // H/HD and D/HD overloading is not allowed. |
| extern "C" __host__ __device__ int chhd1() { return 0; } // expected-note {{previous declaration is here}} |
| extern "C" __host__ int chhd1() { return 0; } |
| // expected-error@-1 {{__host__ function 'chhd1' cannot overload __host__ __device__ function 'chhd1'}} |
| |
| extern "C" __host__ int chhd2() { return 0; } // expected-note {{previous declaration is here}} |
| extern "C" __host__ __device__ int chhd2() { return 0; } |
| // expected-error@-1 {{__host__ __device__ function 'chhd2' cannot overload __host__ function 'chhd2'}} |
| |
| // Helper functions to verify calling restrictions. |
| __device__ DeviceReturnTy d() { return DeviceReturnTy(); } |
| // expected-note@-1 1+ {{'d' declared here}} |
| // expected-note@-2 1+ {{candidate function not viable: call to __device__ function from __host__ function}} |
| // expected-note@-3 0+ {{candidate function not viable: call to __device__ function from __host__ __device__ function}} |
| |
| __host__ HostReturnTy h() { return HostReturnTy(); } |
| // expected-note@-1 1+ {{'h' declared here}} |
| // expected-note@-2 1+ {{candidate function not viable: call to __host__ function from __device__ function}} |
| // expected-note@-3 0+ {{candidate function not viable: call to __host__ function from __host__ __device__ function}} |
| // expected-note@-4 1+ {{candidate function not viable: call to __host__ function from __global__ function}} |
| |
| __global__ void g() {} |
| // expected-note@-1 1+ {{'g' declared here}} |
| // expected-note@-2 1+ {{candidate function not viable: call to __global__ function from __device__ function}} |
| // expected-note@-3 0+ {{candidate function not viable: call to __global__ function from __host__ __device__ function}} |
| // expected-note@-4 1+ {{candidate function not viable: call to __global__ function from __global__ function}} |
| |
| extern "C" __device__ DeviceReturnTy cd() { return DeviceReturnTy(); } |
| // expected-note@-1 1+ {{'cd' declared here}} |
| // expected-note@-2 1+ {{candidate function not viable: call to __device__ function from __host__ function}} |
| // expected-note@-3 0+ {{candidate function not viable: call to __device__ function from __host__ __device__ function}} |
| |
| extern "C" __host__ HostReturnTy ch() { return HostReturnTy(); } |
| // expected-note@-1 1+ {{'ch' declared here}} |
| // expected-note@-2 1+ {{candidate function not viable: call to __host__ function from __device__ function}} |
| // expected-note@-3 0+ {{candidate function not viable: call to __host__ function from __host__ __device__ function}} |
| // expected-note@-4 1+ {{candidate function not viable: call to __host__ function from __global__ function}} |
| |
| __host__ void hostf() { |
| DeviceFnPtr fp_d = d; // expected-error {{reference to __device__ function 'd' in __host__ function}} |
| DeviceReturnTy ret_d = d(); // expected-error {{no matching function for call to 'd'}} |
| DeviceFnPtr fp_cd = cd; // expected-error {{reference to __device__ function 'cd' in __host__ function}} |
| DeviceReturnTy ret_cd = cd(); // expected-error {{no matching function for call to 'cd'}} |
| |
| HostFnPtr fp_h = h; |
| HostReturnTy ret_h = h(); |
| HostFnPtr fp_ch = ch; |
| HostReturnTy ret_ch = ch(); |
| |
| HostFnPtr fp_dh = dh; |
| HostReturnTy ret_dh = dh(); |
| HostFnPtr fp_cdh = cdh; |
| HostReturnTy ret_cdh = cdh(); |
| |
| GlobalFnPtr fp_g = g; |
| g(); // expected-error {{call to global function 'g' not configured}} |
| g<<<0, 0>>>(); |
| } |
| |
| __device__ void devicef() { |
| DeviceFnPtr fp_d = d; |
| DeviceReturnTy ret_d = d(); |
| DeviceFnPtr fp_cd = cd; |
| DeviceReturnTy ret_cd = cd(); |
| |
| HostFnPtr fp_h = h; // expected-error {{reference to __host__ function 'h' in __device__ function}} |
| HostReturnTy ret_h = h(); // expected-error {{no matching function for call to 'h'}} |
| HostFnPtr fp_ch = ch; // expected-error {{reference to __host__ function 'ch' in __device__ function}} |
| HostReturnTy ret_ch = ch(); // expected-error {{no matching function for call to 'ch'}} |
| |
| DeviceFnPtr fp_dh = dh; |
| DeviceReturnTy ret_dh = dh(); |
| DeviceFnPtr fp_cdh = cdh; |
| DeviceReturnTy ret_cdh = cdh(); |
| |
| GlobalFnPtr fp_g = g; // expected-error {{reference to __global__ function 'g' in __device__ function}} |
| g(); // expected-error {{no matching function for call to 'g'}} |
| g<<<0,0>>>(); // expected-error {{reference to __global__ function 'g' in __device__ function}} |
| } |
| |
| __global__ void globalf() { |
| DeviceFnPtr fp_d = d; |
| DeviceReturnTy ret_d = d(); |
| DeviceFnPtr fp_cd = cd; |
| DeviceReturnTy ret_cd = cd(); |
| |
| HostFnPtr fp_h = h; // expected-error {{reference to __host__ function 'h' in __global__ function}} |
| HostReturnTy ret_h = h(); // expected-error {{no matching function for call to 'h'}} |
| HostFnPtr fp_ch = ch; // expected-error {{reference to __host__ function 'ch' in __global__ function}} |
| HostReturnTy ret_ch = ch(); // expected-error {{no matching function for call to 'ch'}} |
| |
| DeviceFnPtr fp_dh = dh; |
| DeviceReturnTy ret_dh = dh(); |
| DeviceFnPtr fp_cdh = cdh; |
| DeviceReturnTy ret_cdh = cdh(); |
| |
| GlobalFnPtr fp_g = g; // expected-error {{reference to __global__ function 'g' in __global__ function}} |
| g(); // expected-error {{no matching function for call to 'g'}} |
| g<<<0,0>>>(); // expected-error {{reference to __global__ function 'g' in __global__ function}} |
| } |
| |
| __host__ __device__ void hostdevicef() { |
| DeviceFnPtr fp_d = d; |
| DeviceReturnTy ret_d = d(); |
| DeviceFnPtr fp_cd = cd; |
| DeviceReturnTy ret_cd = cd(); |
| #if !defined(__CUDA_ARCH__) |
| // expected-error@-5 {{reference to __device__ function 'd' in __host__ __device__ function}} |
| // expected-error@-5 {{reference to __device__ function 'd' in __host__ __device__ function}} |
| // expected-error@-5 {{reference to __device__ function 'cd' in __host__ __device__ function}} |
| // expected-error@-5 {{reference to __device__ function 'cd' in __host__ __device__ function}} |
| #endif |
| |
| HostFnPtr fp_h = h; |
| HostReturnTy ret_h = h(); |
| HostFnPtr fp_ch = ch; |
| HostReturnTy ret_ch = ch(); |
| #if defined(__CUDA_ARCH__) |
| // expected-error@-5 {{reference to __host__ function 'h' in __host__ __device__ function}} |
| // expected-error@-5 {{reference to __host__ function 'h' in __host__ __device__ function}} |
| // expected-error@-5 {{reference to __host__ function 'ch' in __host__ __device__ function}} |
| // expected-error@-5 {{reference to __host__ function 'ch' in __host__ __device__ function}} |
| #endif |
| |
| CurrentFnPtr fp_dh = dh; |
| CurrentReturnTy ret_dh = dh(); |
| CurrentFnPtr fp_cdh = cdh; |
| CurrentReturnTy ret_cdh = cdh(); |
| |
| GlobalFnPtr fp_g = g; |
| #if defined(__CUDA_ARCH__) |
| // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}} |
| #endif |
| |
| g(); |
| #if defined (__CUDA_ARCH__) |
| // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}} |
| #else |
| // expected-error@-4 {{call to global function 'g' not configured}} |
| #endif |
| |
| g<<<0,0>>>(); |
| #if defined(__CUDA_ARCH__) |
| // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}} |
| #endif |
| } |
| |
| // Test for address of overloaded function resolution in the global context. |
| HostFnPtr fp_h = h; |
| HostFnPtr fp_ch = ch; |
| CurrentFnPtr fp_dh = dh; |
| CurrentFnPtr fp_cdh = cdh; |
| GlobalFnPtr fp_g = g; |
| |
| |
| // Test overloading of destructors |
| // Can't mix H and unattributed destructors |
| struct d_h { |
| ~d_h() {} // expected-note {{previous definition is here}} |
| __host__ ~d_h() {} // expected-error {{destructor cannot be redeclared}} |
| }; |
| |
| // HD is OK |
| struct d_hd { |
| __host__ __device__ ~d_hd() {} |
| }; |
| |
| // 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@-1 {{__host__ __device__ function 'operator delete' cannot overload __host__ function 'operator delete'}} |
| }; |
| |
| struct m_hdh { |
| __host__ __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}} |
| __host__ void operator delete(void *ptr) {} |
| // expected-error@-1 {{__host__ function 'operator delete' cannot overload __host__ __device__ function 'operator delete'}} |
| }; |
| |
| struct m_dhd { |
| __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}} |
| __host__ __device__ void operator delete(void *ptr) {} |
| // expected-error@-1 {{__host__ __device__ function 'operator delete' cannot overload __device__ function 'operator delete'}} |
| }; |
| |
| struct m_hdd { |
| __host__ __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}} |
| __device__ void operator delete(void *ptr) {} |
| // expected-error@-1 {{__device__ function 'operator delete' cannot overload __host__ __device__ function 'operator delete'}} |
| }; |
| |
| // __global__ functions can't be overloaded based on attribute |
| // difference. |
| struct G { |
| friend void friend_of_g(G &arg); // expected-note {{previous declaration is here}} |
| private: |
| int x; // expected-note {{declared private here}} |
| }; |
| __global__ void friend_of_g(G &arg) { int x = arg.x; } |
| // expected-error@-1 {{__global__ function 'friend_of_g' cannot overload __host__ function 'friend_of_g'}} |
| // expected-error@-2 {{'x' is a private member of 'G'}} |
| void friend_of_g(G &arg) { int x = arg.x; } |
| |
| // HD functions are sometimes allowed to call H or D functions -- this |
| // is an artifact of the source-to-source splitting performed by nvcc |
| // that we need to mimic. During device mode compilation in nvcc, host |
| // functions aren't present at all, so don't participate in |
| // overloading. But in clang, H and D functions are present in both |
| // compilation modes. Clang normally uses the target attribute as a |
| // tiebreaker between overloads with otherwise identical priority, but |
| // in order to match nvcc's behavior, we sometimes need to wholly |
| // discard overloads that would not be present during compilation |
| // under nvcc. |
| |
| template <typename T> TemplateReturnTy template_vs_function(T arg) { |
| return TemplateReturnTy(); |
| } |
| __device__ DeviceReturnTy template_vs_function(float arg) { |
| return DeviceReturnTy(); |
| } |
| |
| // Here we expect to call the templated function during host compilation, even |
| // if -fcuda-disable-target-call-checks is passed, and even though C++ overload |
| // rules prefer the non-templated function. |
| __host__ __device__ void test_host_device_calls_template(void) { |
| #ifdef __CUDA_ARCH__ |
| typedef DeviceReturnTy ExpectedReturnTy; |
| #else |
| typedef TemplateReturnTy ExpectedReturnTy; |
| #endif |
| |
| ExpectedReturnTy ret1 = template_vs_function(1.0f); |
| ExpectedReturnTy ret2 = template_vs_function(2.0); |
| } |
| |
| // Calls from __host__ and __device__ functions should always call the |
| // overloaded function that matches their mode. |
| __host__ void test_host_calls_template_fn() { |
| TemplateReturnTy ret1 = template_vs_function(1.0f); |
| TemplateReturnTy ret2 = template_vs_function(2.0); |
| } |
| |
| __device__ void test_device_calls_template_fn() { |
| DeviceReturnTy ret1 = template_vs_function(1.0f); |
| DeviceReturnTy ret2 = template_vs_function(2.0); |
| } |
| |
| // If we have a mix of HD and H-only or D-only candidates in the overload set, |
| // normal C++ overload resolution rules apply first. |
| template <typename T> TemplateReturnTy template_vs_hd_function(T arg) |
| #ifdef __CUDA_ARCH__ |
| //expected-note@-2 {{declared here}} |
| #endif |
| { |
| return TemplateReturnTy(); |
| } |
| __host__ __device__ HostDeviceReturnTy template_vs_hd_function(float arg) { |
| return HostDeviceReturnTy(); |
| } |
| |
| __host__ __device__ void test_host_device_calls_hd_template() { |
| HostDeviceReturnTy ret1 = template_vs_hd_function(1.0f); |
| TemplateReturnTy ret2 = template_vs_hd_function(1); |
| #ifdef __CUDA_ARCH__ |
| // expected-error@-2 {{reference to __host__ function 'template_vs_hd_function<int>' in __host__ __device__ function}} |
| #endif |
| } |
| |
| __host__ void test_host_calls_hd_template() { |
| HostDeviceReturnTy ret1 = template_vs_hd_function(1.0f); |
| TemplateReturnTy ret2 = template_vs_hd_function(1); |
| } |
| |
| __device__ void test_device_calls_hd_template() { |
| HostDeviceReturnTy ret1 = template_vs_hd_function(1.0f); |
| // Host-only function template is not callable with strict call checks, |
| // so for device side HD function will be the only choice. |
| HostDeviceReturnTy ret2 = template_vs_hd_function(1); |
| } |
| |
| // Check that overloads still work the same way on both host and |
| // device side when the overload set contains only functions from one |
| // side of compilation. |
| __device__ DeviceReturnTy device_only_function(int arg) { return DeviceReturnTy(); } |
| __device__ DeviceReturnTy2 device_only_function(float arg) { return DeviceReturnTy2(); } |
| #ifndef __CUDA_ARCH__ |
| // expected-note@-3 {{'device_only_function' declared here}} |
| // expected-note@-3 {{'device_only_function' declared here}} |
| #endif |
| __host__ HostReturnTy host_only_function(int arg) { return HostReturnTy(); } |
| __host__ HostReturnTy2 host_only_function(float arg) { return HostReturnTy2(); } |
| #ifdef __CUDA_ARCH__ |
| // expected-note@-3 {{'host_only_function' declared here}} |
| // expected-note@-3 {{'host_only_function' declared here}} |
| #endif |
| |
| __host__ __device__ void test_host_device_single_side_overloading() { |
| DeviceReturnTy ret1 = device_only_function(1); |
| DeviceReturnTy2 ret2 = device_only_function(1.0f); |
| #ifndef __CUDA_ARCH__ |
| // expected-error@-3 {{reference to __device__ function 'device_only_function' in __host__ __device__ function}} |
| // expected-error@-3 {{reference to __device__ function 'device_only_function' in __host__ __device__ function}} |
| #endif |
| HostReturnTy ret3 = host_only_function(1); |
| HostReturnTy2 ret4 = host_only_function(1.0f); |
| #ifdef __CUDA_ARCH__ |
| // expected-error@-3 {{reference to __host__ function 'host_only_function' in __host__ __device__ function}} |
| // expected-error@-3 {{reference to __host__ function 'host_only_function' in __host__ __device__ function}} |
| #endif |
| } |
| |
| // Verify that we allow overloading function templates. |
| template <typename T> __host__ T template_overload(const T &a) { return a; }; |
| template <typename T> __device__ T template_overload(const T &a) { return a; }; |
| |
| __host__ void test_host_template_overload() { |
| template_overload(1); // OK. Attribute-based overloading picks __host__ variant. |
| } |
| __device__ void test_device_template_overload() { |
| template_overload(1); // OK. Attribute-based overloading picks __device__ variant. |
| } |
| |
| // Two classes with `operator-` defined. One of them is device only. |
| struct C1; |
| struct C2; |
| __device__ |
| int operator-(const C1 &x, const C1 &y); |
| int operator-(const C2 &x, const C2 &y); |
| |
| template <typename T> |
| __host__ __device__ int constexpr_overload(const T &x, const T &y) { |
| return x - y; |
| } |
| |
| // Verify that function overloading doesn't prune candidate wrongly. |
| int test_constexpr_overload(C2 &x, C2 &y) { |
| return constexpr_overload(x, y); |
| } |