| // REQUIRES: x86-registered-target |
| // REQUIRES: nvptx-registered-target |
| |
| // RUN: %clang_cc1 -std=c++14 -triple x86_64-unknown-linux-gnu -fsyntax-only \ |
| // RUN: -verify=host,hostdefer,devdefer,expected %s |
| // RUN: %clang_cc1 -std=c++14 -triple nvptx64-nvidia-cuda -fsyntax-only \ |
| // RUN: -fcuda-is-device -verify=dev,devnodeferonly,hostdefer,devdefer,expected %s |
| // RUN: %clang_cc1 -fgpu-exclude-wrong-side-overloads -fgpu-defer-diag -DDEFER=1 \ |
| // RUN: -std=c++14 -triple x86_64-unknown-linux-gnu -fsyntax-only \ |
| // RUN: -verify=host,hostdefer,expected %s |
| // RUN: %clang_cc1 -fgpu-exclude-wrong-side-overloads -fgpu-defer-diag -DDEFER=1 \ |
| // RUN: -std=c++14 -triple nvptx64-nvidia-cuda -fsyntax-only -fcuda-is-device \ |
| // RUN: -verify=dev,devdeferonly,devdefer,expected %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(); } |
| // host-note@-1 1+ {{'d' declared here}} |
| // hostdefer-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(); } |
| // dev-note@-1 1+ {{'h' declared here}} |
| // devdefer-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}} |
| // devdefer-note@-4 1+ {{candidate function not viable: call to __host__ function from __global__ function}} |
| |
| __global__ void g() {} |
| // dev-note@-1 1+ {{'g' declared here}} |
| // devdefer-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}} |
| // devdefer-note@-4 1+ {{candidate function not viable: call to __global__ function from __global__ function}} |
| |
| extern "C" __device__ DeviceReturnTy cd() { return DeviceReturnTy(); } |
| // host-note@-1 1+ {{'cd' declared here}} |
| // hostdefer-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(); } |
| // dev-note@-1 1+ {{'ch' declared here}} |
| // devdefer-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}} |
| // devdefer-note@-4 1+ {{candidate function not viable: call to __host__ function from __global__ function}} |
| |
| __host__ void hostf() { |
| DeviceFnPtr fp_d = d; // host-error {{reference to __device__ function 'd' in __host__ function}} |
| DeviceReturnTy ret_d = d(); // hostdefer-error {{no matching function for call to 'd'}} |
| DeviceFnPtr fp_cd = cd; // host-error {{reference to __device__ function 'cd' in __host__ function}} |
| DeviceReturnTy ret_cd = cd(); // hostdefer-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; // dev-error {{reference to __host__ function 'h' in __device__ function}} |
| HostReturnTy ret_h = h(); // devdefer-error {{no matching function for call to 'h'}} |
| HostFnPtr fp_ch = ch; // dev-error {{reference to __host__ function 'ch' in __device__ function}} |
| HostReturnTy ret_ch = ch(); // devdefer-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; // dev-error {{reference to __global__ function 'g' in __device__ function}} |
| g(); // devdefer-error {{no matching function for call to 'g'}} |
| g<<<0,0>>>(); // dev-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; // dev-error {{reference to __host__ function 'h' in __global__ function}} |
| HostReturnTy ret_h = h(); // devdefer-error {{no matching function for call to 'h'}} |
| HostFnPtr fp_ch = ch; // dev-error {{reference to __host__ function 'ch' in __global__ function}} |
| HostReturnTy ret_ch = ch(); // devdefer-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; // dev-error {{reference to __global__ function 'g' in __global__ function}} |
| g(); // devdefer-error {{no matching function for call to 'g'}} |
| g<<<0,0>>>(); // dev-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}} |
| // devdefer-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) |
| // devnodeferonly-note@-1{{'template_vs_hd_function<int>' declared here}} |
| { |
| return TemplateReturnTy(); |
| } |
| __host__ __device__ HostDeviceReturnTy template_vs_hd_function(float arg) { |
| return HostDeviceReturnTy(); |
| } |
| |
| __host__ __device__ void test_host_device_calls_hd_template() { |
| #if __CUDA_ARCH__ && DEFER |
| typedef HostDeviceReturnTy ExpectedReturnTy; |
| #else |
| typedef TemplateReturnTy ExpectedReturnTy; |
| #endif |
| HostDeviceReturnTy ret1 = template_vs_hd_function(1.0f); |
| ExpectedReturnTy ret2 = template_vs_hd_function(1); |
| // devnodeferonly-error@-1{{reference to __host__ function 'template_vs_hd_function<int>' in __host__ __device__ function}} |
| } |
| |
| __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 2{{'device_only_function' declared here}} |
| // expected-note@-3 2{{'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 2{{'host_only_function' declared here}} |
| // expected-note@-3 2{{'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 |
| } |
| |
| // wrong-sided overloading should not cause diagnostic unless it is emitted. |
| // This inline function is not emitted. |
| inline __host__ __device__ void test_host_device_wrong_side_overloading_inline_no_diag() { |
| DeviceReturnTy ret1 = device_only_function(1); |
| DeviceReturnTy2 ret2 = device_only_function(1.0f); |
| HostReturnTy ret3 = host_only_function(1); |
| HostReturnTy2 ret4 = host_only_function(1.0f); |
| } |
| |
| // wrong-sided overloading should cause diagnostic if it is emitted. |
| // This inline function is emitted since it is called by an emitted function. |
| inline __host__ __device__ void test_host_device_wrong_side_overloading_inline_diag() { |
| 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 |
| } |
| |
| __host__ __device__ void test_host_device_wrong_side_overloading_inline_diag_caller() { |
| test_host_device_wrong_side_overloading_inline_diag(); |
| // expected-note@-1 {{called by 'test_host_device_wrong_side_overloading_inline_diag_caller'}} |
| } |
| |
| // 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); |
| } |
| |
| // Verify no ambiguity for new operator. |
| void *a = new int; |
| __device__ void *b = new int; |
| // expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables.}} |
| |
| // Verify no ambiguity for new operator. |
| template<typename _Tp> _Tp&& f(); |
| template<typename _Tp, typename = decltype(new _Tp(f<_Tp>()))> |
| void __test(); |
| |
| void foo() { |
| __test<int>(); |
| } |
| |
| // Test resolving implicit host device candidate vs wrong-sided candidate. |
| // In device compilation, implicit host device caller choose implicit host |
| // device candidate and wrong-sided candidate with equal preference. |
| // Resolution result should not change with/without pragma. |
| namespace ImplicitHostDeviceVsWrongSided { |
| HostReturnTy callee(double x); |
| #pragma clang force_cuda_host_device begin |
| HostDeviceReturnTy callee(int x); |
| inline HostReturnTy implicit_hd_caller() { |
| return callee(1.0); |
| } |
| #pragma clang force_cuda_host_device end |
| } |
| |
| // Test resolving implicit host device candidate vs same-sided candidate. |
| // In host compilation, implicit host device caller choose implicit host |
| // device candidate and same-sided candidate with equal preference. |
| // Resolution result should not change with/without pragma. |
| namespace ImplicitHostDeviceVsSameSide { |
| HostReturnTy callee(int x); |
| #pragma clang force_cuda_host_device begin |
| HostDeviceReturnTy callee(double x); |
| inline HostDeviceReturnTy implicit_hd_caller() { |
| return callee(1.0); |
| } |
| #pragma clang force_cuda_host_device end |
| } |
| |
| // Test resolving explicit host device candidate vs. wrong-sided candidate. |
| // When -fgpu-defer-diag is off, wrong-sided candidate is not excluded, therefore |
| // the first callee is chosen. |
| // When -fgpu-defer-diag is on, wrong-sided candidate is excluded, therefore |
| // the second callee is chosen. |
| namespace ExplicitHostDeviceVsWrongSided { |
| HostReturnTy callee(double x); |
| __host__ __device__ HostDeviceReturnTy callee(int x); |
| #if __CUDA_ARCH__ && DEFER |
| typedef HostDeviceReturnTy ExpectedRetTy; |
| #else |
| typedef HostReturnTy ExpectedRetTy; |
| #endif |
| inline __host__ __device__ ExpectedRetTy explicit_hd_caller() { |
| return callee(1.0); |
| } |
| } |
| |
| // In the implicit host device function 'caller', the second 'callee' should be |
| // chosen since it has better match, even though it is an implicit host device |
| // function whereas the first 'callee' is a host function. A diagnostic will be |
| // emitted if the first 'callee' is chosen since deduced return type cannot be |
| // used before it is defined. |
| namespace ImplicitHostDeviceByConstExpr { |
| template <class a> a b; |
| auto callee(...); |
| template <class d> constexpr auto callee(d) -> decltype(0); |
| struct e { |
| template <class ad, class... f> static auto g(ad, f...) { |
| return h<e, decltype(b<f>)...>; |
| } |
| struct i { |
| template <class, class... f> static constexpr auto caller(f... k) { |
| return callee(k...); |
| } |
| }; |
| template <class, class... f> static auto h() { |
| return i::caller<int, f...>; |
| } |
| }; |
| class l { |
| l() { |
| e::g([] {}, this); |
| } |
| }; |
| } |
| |
| // Implicit HD candidate competes with device candidate. |
| // a and b have implicit HD copy ctor. In copy ctor of b, ctor of a is resolved. |
| // copy ctor of a should win over a(short), otherwise there will be ambiguity |
| // due to conversion operator. |
| namespace TestImplicitHDWithD { |
| struct a { |
| __device__ a(short); |
| __device__ operator unsigned() const; |
| __device__ operator int() const; |
| }; |
| struct b { |
| a d; |
| }; |
| void f(b g) { b e = g; } |
| } |
| |
| // Implicit HD candidate competes with host candidate. |
| // a and b have implicit HD copy ctor. In copy ctor of b, ctor of a is resolved. |
| // copy ctor of a should win over a(short), otherwise there will be ambiguity |
| // due to conversion operator. |
| namespace TestImplicitHDWithH { |
| struct a { |
| a(short); |
| __device__ operator unsigned() const; |
| __device__ operator int() const; |
| }; |
| struct b { |
| a d; |
| }; |
| void f(b g) { b e = g; } |
| } |
| |
| // Implicit HD candidate competes with HD candidate. |
| // a and b have implicit HD copy ctor. In copy ctor of b, ctor of a is resolved. |
| // copy ctor of a should win over a(short), otherwise there will be ambiguity |
| // due to conversion operator. |
| namespace TestImplicitHDWithHD { |
| struct a { |
| __host__ __device__ a(short); |
| __device__ operator unsigned() const; |
| __device__ operator int() const; |
| }; |
| struct b { |
| a d; |
| }; |
| void f(b g) { b e = g; } |
| } |
| |
| // HD candidate competes with H candidate. |
| // HD has type mismatch whereas H has type match. |
| // In device compilation, H wins when -fgpu-defer-diag is off and HD wins |
| // when -fgpu-defer-diags is on. In both cases the diagnostic should be |
| // deferred. |
| namespace TestDeferNoMatchingFuncNotEmitted { |
| template <typename> struct a {}; |
| namespace b { |
| struct c : a<int> {}; |
| template <typename d> void ag(d); |
| } // namespace b |
| template <typename ae> |
| __host__ __device__ void ag(a<ae>) { |
| ae e; |
| ag(e); |
| } |
| void f() { (void)ag<b::c>; } |
| } |
| |
| namespace TestDeferNoMatchingFuncEmitted { |
| template <typename> struct a {}; |
| namespace b { |
| struct c : a<int> {}; |
| template <typename d> void ag(d); |
| // devnodeferonly-note@-1{{'ag<TestDeferNoMatchingFuncEmitted::b::c>' declared here}} |
| } // namespace b |
| template <typename ae> |
| __host__ __device__ void ag(a<ae>) { |
| ae e; |
| ag(e); |
| // devnodeferonly-error@-1{{reference to __host__ function 'ag<TestDeferNoMatchingFuncEmitted::b::c>' in __host__ __device__ function}} |
| // devdeferonly-error@-2{{no matching function for call to 'ag'}} |
| // devdeferonly-note@-3{{called by 'ag<TestDeferNoMatchingFuncEmitted::b::c>'}} |
| } |
| __host__ __device__ void f() { (void)ag<b::c>; } |
| // devnodeferonly-note@-1{{called by 'f'}} |
| // devdeferonly-note@-2{{called by 'f'}} |
| } |
| |
| // Two HD candidates compete with H candidate. |
| // HDs have type mismatch whereas H has type match. |
| // In device compilation, H wins when -fgpu-defer-diag is off and two HD win |
| // when -fgpu-defer-diags is on. In both cases the diagnostic should be |
| // deferred. |
| namespace TestDeferAmbiguityNotEmitted { |
| template <typename> struct a {}; |
| namespace b { |
| struct c : a<int> {}; |
| template <typename d> void ag(d, int); |
| } // namespace b |
| template <typename ae> |
| __host__ __device__ void ag(a<ae>, float) { |
| ae e; |
| ag(e, 1); |
| } |
| template <typename ae> |
| __host__ __device__ void ag(a<ae>, double) { |
| } |
| void f() { |
| b::c x; |
| ag(x, 1); |
| } |
| } |
| |
| namespace TestDeferAmbiguityEmitted { |
| template <typename> struct a {}; |
| namespace b { |
| struct c : a<int> {}; |
| template <typename d> void ag(d, int); |
| // devnodeferonly-note@-1{{'ag<TestDeferAmbiguityEmitted::b::c>' declared here}} |
| } // namespace b |
| template <typename ae> |
| __host__ __device__ void ag(a<ae>, float) { |
| // devdeferonly-note@-1{{candidate function [with ae = int]}} |
| ae e; |
| ag(e, 1); |
| } |
| template <typename ae> |
| __host__ __device__ void ag(a<ae>, double) { |
| // devdeferonly-note@-1{{candidate function [with ae = int]}} |
| } |
| __host__ __device__ void f() { |
| b::c x; |
| ag(x, 1); |
| // devnodeferonly-error@-1{{reference to __host__ function 'ag<TestDeferAmbiguityEmitted::b::c>' in __host__ __device__ function}} |
| // devdeferonly-error@-2{{call to 'ag' is ambiguous}} |
| } |
| } |
| |
| // Implicit HD functions compute with H function and D function. |
| // In host compilation, foo(0.0, 2) should resolve to X::foo<double, int>. |
| // In device compilation, foo(0.0, 2) should resolve to foo(double, int). |
| // In either case there should be no ambiguity. |
| namespace TestImplicitHDWithHAndD { |
| namespace X { |
| inline double foo(double, double) { return 0;} |
| inline constexpr float foo(float, float) { return 1;} |
| inline constexpr long double foo(long double, long double) { return 2;} |
| template<typename _Tp, typename _Up> inline constexpr double foo(_Tp, _Up) { return 3;} |
| }; |
| using X::foo; |
| inline __device__ double foo(double, double) { return 4;} |
| inline __device__ float foo(float, int) { return 5;} |
| inline __device__ float foo(int, int) { return 6;} |
| inline __device__ double foo(double, int) { return 7;} |
| inline __device__ float foo(float, float) { return 9;} |
| template<typename _Tp, typename _Up> inline __device__ double foo(_Tp, _Up) { return 10;} |
| |
| int g() { |
| return [](){ |
| return foo(0.0, 2); |
| }(); |
| } |
| } |