714 lines
28 KiB
Plaintext
714 lines
28 KiB
Plaintext
// 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);
|
|
}();
|
|
}
|
|
}
|