707 lines · plain
1// RUN: %clang_cc1 -std=c++14 -triple x86_64-unknown-linux-gnu -fsyntax-only \2// RUN: -verify=host,hostdefer,devdefer,expected %s3// RUN: %clang_cc1 -std=c++14 -triple nvptx64-nvidia-cuda -fsyntax-only \4// RUN: -fcuda-is-device -verify=dev,devnodeferonly,hostdefer,devdefer,expected %s5// RUN: %clang_cc1 -fgpu-exclude-wrong-side-overloads -fgpu-defer-diag -DDEFER=1 \6// RUN: -std=c++14 -triple x86_64-unknown-linux-gnu -fsyntax-only \7// RUN: -verify=host,hostdefer,expected %s8// RUN: %clang_cc1 -fgpu-exclude-wrong-side-overloads -fgpu-defer-diag -DDEFER=1 \9// RUN: -std=c++14 -triple nvptx64-nvidia-cuda -fsyntax-only -fcuda-is-device \10// RUN: -verify=dev,devdeferonly,devdefer,expected %s11 12#include "Inputs/cuda.h"13 14// Opaque return types used to check that we pick the right overloads.15struct HostReturnTy {};16struct HostReturnTy2 {};17struct DeviceReturnTy {};18struct DeviceReturnTy2 {};19struct HostDeviceReturnTy {};20struct TemplateReturnTy {};21 22typedef HostReturnTy (*HostFnPtr)();23typedef DeviceReturnTy (*DeviceFnPtr)();24typedef HostDeviceReturnTy (*HostDeviceFnPtr)();25typedef void (*GlobalFnPtr)(); // __global__ functions must return void.26 27// CurrentReturnTy is {HostReturnTy,DeviceReturnTy} during {host,device}28// compilation.29#ifdef __CUDA_ARCH__30typedef DeviceReturnTy CurrentReturnTy;31#else32typedef HostReturnTy CurrentReturnTy;33#endif34 35// CurrentFnPtr is a function pointer to a {host,device} function during36// {host,device} compilation.37typedef CurrentReturnTy (*CurrentFnPtr)();38 39// Host and unattributed functions can't be overloaded.40__host__ void hh() {} // expected-note {{previous definition is here}}41void hh() {} // expected-error {{redefinition of 'hh'}}42 43// H/D overloading is OK.44__host__ HostReturnTy dh() { return HostReturnTy(); }45__device__ DeviceReturnTy dh() { return DeviceReturnTy(); }46 47// H/HD and D/HD are not allowed.48__host__ __device__ int hdh() { return 0; } // expected-note {{previous declaration is here}}49__host__ int hdh() { return 0; }50// expected-error@-1 {{__host__ function 'hdh' cannot overload __host__ __device__ function 'hdh'}}51 52__host__ int hhd() { return 0; } // expected-note {{previous declaration is here}}53__host__ __device__ int hhd() { return 0; }54// expected-error@-1 {{__host__ __device__ function 'hhd' cannot overload __host__ function 'hhd'}}55 56__host__ __device__ int hdd() { return 0; } // expected-note {{previous declaration is here}}57__device__ int hdd() { return 0; }58// expected-error@-1 {{__device__ function 'hdd' cannot overload __host__ __device__ function 'hdd'}}59 60__device__ int dhd() { return 0; } // expected-note {{previous declaration is here}}61__host__ __device__ int dhd() { return 0; }62// expected-error@-1 {{__host__ __device__ function 'dhd' cannot overload __device__ function 'dhd'}}63 64// Same tests for extern "C" functions.65extern "C" __host__ int chh() { return 0; } // expected-note {{previous definition is here}}66extern "C" int chh() { return 0; } // expected-error {{redefinition of 'chh'}}67 68// H/D overloading is OK.69extern "C" __device__ DeviceReturnTy cdh() { return DeviceReturnTy(); }70extern "C" __host__ HostReturnTy cdh() { return HostReturnTy(); }71 72// H/HD and D/HD overloading is not allowed.73extern "C" __host__ __device__ int chhd1() { return 0; } // expected-note {{previous declaration is here}}74extern "C" __host__ int chhd1() { return 0; }75// expected-error@-1 {{__host__ function 'chhd1' cannot overload __host__ __device__ function 'chhd1'}}76 77extern "C" __host__ int chhd2() { return 0; } // expected-note {{previous declaration is here}}78extern "C" __host__ __device__ int chhd2() { return 0; }79// expected-error@-1 {{__host__ __device__ function 'chhd2' cannot overload __host__ function 'chhd2'}}80 81// Helper functions to verify calling restrictions.82__device__ DeviceReturnTy d() { return DeviceReturnTy(); }83// host-note@-1 1+ {{'d' declared here}}84// hostdefer-note@-2 1+ {{candidate function not viable: call to __device__ function from __host__ function}}85// expected-note@-3 0+ {{candidate function not viable: call to __device__ function from __host__ __device__ function}}86 87__host__ HostReturnTy h() { return HostReturnTy(); }88// dev-note@-1 1+ {{'h' declared here}}89// devdefer-note@-2 1+ {{candidate function not viable: call to __host__ function from __device__ function}}90// expected-note@-3 0+ {{candidate function not viable: call to __host__ function from __host__ __device__ function}}91// devdefer-note@-4 1+ {{candidate function not viable: call to __host__ function from __global__ function}}92 93__global__ void g() {}94// expected-note@-3 0+ {{candidate function not viable: call to __global__ function from __host__ __device__ function}}95 96extern "C" __device__ DeviceReturnTy cd() { return DeviceReturnTy(); }97// host-note@-1 1+ {{'cd' declared here}}98// hostdefer-note@-2 1+ {{candidate function not viable: call to __device__ function from __host__ function}}99// expected-note@-3 0+ {{candidate function not viable: call to __device__ function from __host__ __device__ function}}100 101extern "C" __host__ HostReturnTy ch() { return HostReturnTy(); }102// dev-note@-1 1+ {{'ch' declared here}}103// devdefer-note@-2 1+ {{candidate function not viable: call to __host__ function from __device__ function}}104// expected-note@-3 0+ {{candidate function not viable: call to __host__ function from __host__ __device__ function}}105// devdefer-note@-4 1+ {{candidate function not viable: call to __host__ function from __global__ function}}106 107__host__ void hostf() {108 DeviceFnPtr fp_d = d; // host-error {{reference to __device__ function 'd' in __host__ function}}109 DeviceReturnTy ret_d = d(); // hostdefer-error {{no matching function for call to 'd'}}110 DeviceFnPtr fp_cd = cd; // host-error {{reference to __device__ function 'cd' in __host__ function}}111 DeviceReturnTy ret_cd = cd(); // hostdefer-error {{no matching function for call to 'cd'}}112 113 HostFnPtr fp_h = h;114 HostReturnTy ret_h = h();115 HostFnPtr fp_ch = ch;116 HostReturnTy ret_ch = ch();117 118 HostFnPtr fp_dh = dh;119 HostReturnTy ret_dh = dh();120 HostFnPtr fp_cdh = cdh;121 HostReturnTy ret_cdh = cdh();122 123 GlobalFnPtr fp_g = g;124 g(); // expected-error {{call to global function 'g' not configured}}125 g<<<0, 0>>>();126}127 128__device__ void devicef() {129 DeviceFnPtr fp_d = d;130 DeviceReturnTy ret_d = d();131 DeviceFnPtr fp_cd = cd;132 DeviceReturnTy ret_cd = cd();133 134 HostFnPtr fp_h = h; // dev-error {{reference to __host__ function 'h' in __device__ function}}135 HostReturnTy ret_h = h(); // devdefer-error {{no matching function for call to 'h'}}136 HostFnPtr fp_ch = ch; // dev-error {{reference to __host__ function 'ch' in __device__ function}}137 HostReturnTy ret_ch = ch(); // devdefer-error {{no matching function for call to 'ch'}}138 139 DeviceFnPtr fp_dh = dh;140 DeviceReturnTy ret_dh = dh();141 DeviceFnPtr fp_cdh = cdh;142 DeviceReturnTy ret_cdh = cdh();143 144 GlobalFnPtr fp_g = g;145 g(); // expected-error {{call to global function 'g' not configured}}146 g<<<0,0>>>(); // expected-error {{kernel launch from __device__ or __global__ function requires relocatable device code (i.e. requires -fgpu-rdc)}}147}148 149__global__ void globalf() {150 DeviceFnPtr fp_d = d;151 DeviceReturnTy ret_d = d();152 DeviceFnPtr fp_cd = cd;153 DeviceReturnTy ret_cd = cd();154 155 HostFnPtr fp_h = h; // dev-error {{reference to __host__ function 'h' in __global__ function}}156 HostReturnTy ret_h = h(); // devdefer-error {{no matching function for call to 'h'}}157 HostFnPtr fp_ch = ch; // dev-error {{reference to __host__ function 'ch' in __global__ function}}158 HostReturnTy ret_ch = ch(); // devdefer-error {{no matching function for call to 'ch'}}159 160 DeviceFnPtr fp_dh = dh;161 DeviceReturnTy ret_dh = dh();162 DeviceFnPtr fp_cdh = cdh;163 DeviceReturnTy ret_cdh = cdh();164 165 GlobalFnPtr fp_g = g;166 g(); // expected-error {{call to global function 'g' not configured}}167 g<<<0,0>>>(); // expected-error {{kernel launch from __device__ or __global__ function requires relocatable device code (i.e. requires -fgpu-rdc)}}168}169 170__host__ __device__ void hostdevicef() {171 DeviceFnPtr fp_d = d;172 DeviceReturnTy ret_d = d();173 DeviceFnPtr fp_cd = cd;174 DeviceReturnTy ret_cd = cd();175#if !defined(__CUDA_ARCH__)176 // expected-error@-5 {{reference to __device__ function 'd' in __host__ __device__ function}}177 // expected-error@-5 {{reference to __device__ function 'd' in __host__ __device__ function}}178 // expected-error@-5 {{reference to __device__ function 'cd' in __host__ __device__ function}}179 // expected-error@-5 {{reference to __device__ function 'cd' in __host__ __device__ function}}180#endif181 182 HostFnPtr fp_h = h;183 HostReturnTy ret_h = h();184 HostFnPtr fp_ch = ch;185 HostReturnTy ret_ch = ch();186#if defined(__CUDA_ARCH__)187 // expected-error@-5 {{reference to __host__ function 'h' in __host__ __device__ function}}188 // expected-error@-5 {{reference to __host__ function 'h' in __host__ __device__ function}}189 // devdefer-error@-5 {{reference to __host__ function 'ch' in __host__ __device__ function}}190 // expected-error@-5 {{reference to __host__ function 'ch' in __host__ __device__ function}}191#endif192 193 CurrentFnPtr fp_dh = dh;194 CurrentReturnTy ret_dh = dh();195 CurrentFnPtr fp_cdh = cdh;196 CurrentReturnTy ret_cdh = cdh();197 198 GlobalFnPtr fp_g = g;199 200 g();201 // expected-error@-1 {{call to global function 'g' not configured}}202 203 g<<<0,0>>>();204#if defined(__CUDA_ARCH__)205 // expected-error@-2 {{kernel launch from __device__ or __global__ function requires relocatable device code (i.e. requires -fgpu-rdc)}}206#endif207}208 209// Test for address of overloaded function resolution in the global context.210HostFnPtr fp_h = h;211HostFnPtr fp_ch = ch;212#if defined (__CUDA_ARCH__)213__device__214#endif215CurrentFnPtr fp_dh = dh;216#if defined (__CUDA_ARCH__)217__device__218#endif219CurrentFnPtr fp_cdh = cdh;220GlobalFnPtr fp_g = g;221 222 223// Test overloading of destructors224// Can't mix H and unattributed destructors225struct d_h {226 ~d_h() {} // expected-note {{previous definition is here}}227 __host__ ~d_h() {} // expected-error {{destructor cannot be redeclared}}228};229 230// HD is OK231struct d_hd {232 __host__ __device__ ~d_hd() {}233};234 235// Test overloading of member functions236struct m_h {237 void operator delete(void *ptr); // expected-note {{previous declaration is here}}238 __host__ void operator delete(void *ptr); // expected-error {{class member cannot be redeclared}}239};240 241// D/H overloading is OK242struct m_dh {243 __device__ void operator delete(void *ptr);244 __host__ void operator delete(void *ptr);245};246 247// HD by itself is OK248struct m_hd {249 __device__ __host__ void operator delete(void *ptr);250};251 252struct m_hhd {253 __host__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}254 __host__ __device__ void operator delete(void *ptr) {}255 // expected-error@-1 {{__host__ __device__ function 'operator delete' cannot overload __host__ function 'operator delete'}}256};257 258struct m_hdh {259 __host__ __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}260 __host__ void operator delete(void *ptr) {}261 // expected-error@-1 {{__host__ function 'operator delete' cannot overload __host__ __device__ function 'operator delete'}}262};263 264struct m_dhd {265 __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}266 __host__ __device__ void operator delete(void *ptr) {}267 // expected-error@-1 {{__host__ __device__ function 'operator delete' cannot overload __device__ function 'operator delete'}}268};269 270struct m_hdd {271 __host__ __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}272 __device__ void operator delete(void *ptr) {}273 // expected-error@-1 {{__device__ function 'operator delete' cannot overload __host__ __device__ function 'operator delete'}}274};275 276// __global__ functions can't be overloaded based on attribute277// difference.278struct G {279 friend void friend_of_g(G &arg); // expected-note {{previous declaration is here}}280private:281 int x; // expected-note {{declared private here}}282};283__global__ void friend_of_g(G &arg) { int x = arg.x; }284// expected-error@-1 {{__global__ function 'friend_of_g' cannot overload __host__ function 'friend_of_g'}}285// expected-error@-2 {{'x' is a private member of 'G'}}286void friend_of_g(G &arg) { int x = arg.x; }287 288// HD functions are sometimes allowed to call H or D functions -- this289// is an artifact of the source-to-source splitting performed by nvcc290// that we need to mimic. During device mode compilation in nvcc, host291// functions aren't present at all, so don't participate in292// overloading. But in clang, H and D functions are present in both293// compilation modes. Clang normally uses the target attribute as a294// tiebreaker between overloads with otherwise identical priority, but295// in order to match nvcc's behavior, we sometimes need to wholly296// discard overloads that would not be present during compilation297// under nvcc.298 299template <typename T> TemplateReturnTy template_vs_function(T arg) {300 return TemplateReturnTy();301}302__device__ DeviceReturnTy template_vs_function(float arg) {303 return DeviceReturnTy();304}305 306// Here we expect to call the templated function during host compilation, even307// if -fcuda-disable-target-call-checks is passed, and even though C++ overload308// rules prefer the non-templated function.309__host__ __device__ void test_host_device_calls_template(void) {310#ifdef __CUDA_ARCH__311 typedef DeviceReturnTy ExpectedReturnTy;312#else313 typedef TemplateReturnTy ExpectedReturnTy;314#endif315 316 ExpectedReturnTy ret1 = template_vs_function(1.0f);317 ExpectedReturnTy ret2 = template_vs_function(2.0);318}319 320// Calls from __host__ and __device__ functions should always call the321// overloaded function that matches their mode.322__host__ void test_host_calls_template_fn() {323 TemplateReturnTy ret1 = template_vs_function(1.0f);324 TemplateReturnTy ret2 = template_vs_function(2.0);325}326 327__device__ void test_device_calls_template_fn() {328 DeviceReturnTy ret1 = template_vs_function(1.0f);329 DeviceReturnTy ret2 = template_vs_function(2.0);330}331 332// If we have a mix of HD and H-only or D-only candidates in the overload set,333// normal C++ overload resolution rules apply first.334template <typename T> TemplateReturnTy template_vs_hd_function(T arg)335// devnodeferonly-note@-1{{'template_vs_hd_function<int>' declared here}}336{337 return TemplateReturnTy();338}339__host__ __device__ HostDeviceReturnTy template_vs_hd_function(float arg) {340 return HostDeviceReturnTy();341}342 343__host__ __device__ void test_host_device_calls_hd_template() {344#if __CUDA_ARCH__ && DEFER345 typedef HostDeviceReturnTy ExpectedReturnTy;346#else347 typedef TemplateReturnTy ExpectedReturnTy;348#endif349 HostDeviceReturnTy ret1 = template_vs_hd_function(1.0f);350 ExpectedReturnTy ret2 = template_vs_hd_function(1);351 // devnodeferonly-error@-1{{reference to __host__ function 'template_vs_hd_function<int>' in __host__ __device__ function}}352}353 354__host__ void test_host_calls_hd_template() {355 HostDeviceReturnTy ret1 = template_vs_hd_function(1.0f);356 TemplateReturnTy ret2 = template_vs_hd_function(1);357}358 359__device__ void test_device_calls_hd_template() {360 HostDeviceReturnTy ret1 = template_vs_hd_function(1.0f);361 // Host-only function template is not callable with strict call checks,362 // so for device side HD function will be the only choice.363 HostDeviceReturnTy ret2 = template_vs_hd_function(1);364}365 366// Check that overloads still work the same way on both host and367// device side when the overload set contains only functions from one368// side of compilation.369__device__ DeviceReturnTy device_only_function(int arg) { return DeviceReturnTy(); }370__device__ DeviceReturnTy2 device_only_function(float arg) { return DeviceReturnTy2(); }371#ifndef __CUDA_ARCH__372 // expected-note@-3 2{{'device_only_function' declared here}}373 // expected-note@-3 2{{'device_only_function' declared here}}374#endif375__host__ HostReturnTy host_only_function(int arg) { return HostReturnTy(); }376__host__ HostReturnTy2 host_only_function(float arg) { return HostReturnTy2(); }377#ifdef __CUDA_ARCH__378 // expected-note@-3 2{{'host_only_function' declared here}}379 // expected-note@-3 2{{'host_only_function' declared here}}380#endif381 382__host__ __device__ void test_host_device_single_side_overloading() {383 DeviceReturnTy ret1 = device_only_function(1);384 DeviceReturnTy2 ret2 = device_only_function(1.0f);385#ifndef __CUDA_ARCH__386 // expected-error@-3 {{reference to __device__ function 'device_only_function' in __host__ __device__ function}}387 // expected-error@-3 {{reference to __device__ function 'device_only_function' in __host__ __device__ function}}388#endif389 HostReturnTy ret3 = host_only_function(1);390 HostReturnTy2 ret4 = host_only_function(1.0f);391#ifdef __CUDA_ARCH__392 // expected-error@-3 {{reference to __host__ function 'host_only_function' in __host__ __device__ function}}393 // expected-error@-3 {{reference to __host__ function 'host_only_function' in __host__ __device__ function}}394#endif395}396 397// wrong-sided overloading should not cause diagnostic unless it is emitted.398// This inline function is not emitted.399inline __host__ __device__ void test_host_device_wrong_side_overloading_inline_no_diag() {400 DeviceReturnTy ret1 = device_only_function(1);401 DeviceReturnTy2 ret2 = device_only_function(1.0f);402 HostReturnTy ret3 = host_only_function(1);403 HostReturnTy2 ret4 = host_only_function(1.0f);404}405 406// wrong-sided overloading should cause diagnostic if it is emitted.407// This inline function is emitted since it is called by an emitted function.408inline __host__ __device__ void test_host_device_wrong_side_overloading_inline_diag() {409 DeviceReturnTy ret1 = device_only_function(1);410 DeviceReturnTy2 ret2 = device_only_function(1.0f);411#ifndef __CUDA_ARCH__412 // expected-error@-3 {{reference to __device__ function 'device_only_function' in __host__ __device__ function}}413 // expected-error@-3 {{reference to __device__ function 'device_only_function' in __host__ __device__ function}}414#endif415 HostReturnTy ret3 = host_only_function(1);416 HostReturnTy2 ret4 = host_only_function(1.0f);417#ifdef __CUDA_ARCH__418 // expected-error@-3 {{reference to __host__ function 'host_only_function' in __host__ __device__ function}}419 // expected-error@-3 {{reference to __host__ function 'host_only_function' in __host__ __device__ function}}420#endif421}422 423__host__ __device__ void test_host_device_wrong_side_overloading_inline_diag_caller() {424 test_host_device_wrong_side_overloading_inline_diag();425 // expected-note@-1 {{called by 'test_host_device_wrong_side_overloading_inline_diag_caller'}}426}427 428// Verify that we allow overloading function templates.429template <typename T> __host__ T template_overload(const T &a) { return a; };430template <typename T> __device__ T template_overload(const T &a) { return a; };431 432__host__ void test_host_template_overload() {433 template_overload(1); // OK. Attribute-based overloading picks __host__ variant.434}435__device__ void test_device_template_overload() {436 template_overload(1); // OK. Attribute-based overloading picks __device__ variant.437}438 439// Two classes with `operator-` defined. One of them is device only.440struct C1;441struct C2;442__device__443int operator-(const C1 &x, const C1 &y);444int operator-(const C2 &x, const C2 &y);445 446template <typename T>447__host__ __device__ int constexpr_overload(const T &x, const T &y) {448 return x - y;449}450 451// Verify that function overloading doesn't prune candidate wrongly.452int test_constexpr_overload(C2 &x, C2 &y) {453 return constexpr_overload(x, y);454}455 456// Verify no ambiguity for new operator.457void *a = new int;458__device__ void *b = new int;459// expected-error@-1{{dynamic initialization is not supported for __device__, __constant__, __shared__, and __managed__ variables}}460 461// Verify no ambiguity for new operator.462template<typename _Tp> _Tp&& f();463template<typename _Tp, typename = decltype(new _Tp(f<_Tp>()))>464void __test();465 466void foo() {467 __test<int>();468}469 470// Test resolving implicit host device candidate vs wrong-sided candidate.471// In device compilation, implicit host device caller choose implicit host472// device candidate and wrong-sided candidate with equal preference.473// Resolution result should not change with/without pragma.474namespace ImplicitHostDeviceVsWrongSided {475HostReturnTy callee(double x);476#pragma clang force_cuda_host_device begin477HostDeviceReturnTy callee(int x);478inline HostReturnTy implicit_hd_caller() {479 return callee(1.0);480}481#pragma clang force_cuda_host_device end482}483 484// Test resolving implicit host device candidate vs same-sided candidate.485// In host compilation, implicit host device caller choose implicit host486// device candidate and same-sided candidate with equal preference.487// Resolution result should not change with/without pragma.488namespace ImplicitHostDeviceVsSameSide {489HostReturnTy callee(int x);490#pragma clang force_cuda_host_device begin491HostDeviceReturnTy callee(double x);492inline HostDeviceReturnTy implicit_hd_caller() {493 return callee(1.0);494}495#pragma clang force_cuda_host_device end496}497 498// Test resolving explicit host device candidate vs. wrong-sided candidate.499// When -fgpu-defer-diag is off, wrong-sided candidate is not excluded, therefore500// the first callee is chosen.501// When -fgpu-defer-diag is on, wrong-sided candidate is excluded, therefore502// the second callee is chosen.503namespace ExplicitHostDeviceVsWrongSided {504HostReturnTy callee(double x);505__host__ __device__ HostDeviceReturnTy callee(int x);506#if __CUDA_ARCH__ && DEFER507typedef HostDeviceReturnTy ExpectedRetTy;508#else509typedef HostReturnTy ExpectedRetTy;510#endif511inline __host__ __device__ ExpectedRetTy explicit_hd_caller() {512 return callee(1.0);513}514}515 516// In the implicit host device function 'caller', the second 'callee' should be517// chosen since it has better match, even though it is an implicit host device518// function whereas the first 'callee' is a host function. A diagnostic will be519// emitted if the first 'callee' is chosen since deduced return type cannot be520// used before it is defined.521namespace ImplicitHostDeviceByConstExpr {522template <class a> a b;523auto callee(...);524template <class d> constexpr auto callee(d) -> decltype(0);525struct e {526 template <class ad, class... f> static auto g(ad, f...) {527 return h<e, decltype(b<f>)...>;528 }529 struct i {530 template <class, class... f> static constexpr auto caller(f... k) {531 return callee(k...);532 }533 };534 template <class, class... f> static auto h() {535 return i::caller<int, f...>;536 }537};538class l {539 l() {540 e::g([] {}, this);541 }542};543}544 545// Implicit HD candidate competes with device candidate.546// a and b have implicit HD copy ctor. In copy ctor of b, ctor of a is resolved.547// copy ctor of a should win over a(short), otherwise there will be ambiguity548// due to conversion operator.549namespace TestImplicitHDWithD {550 struct a {551 __device__ a(short);552 __device__ operator unsigned() const;553 __device__ operator int() const;554 };555 struct b {556 a d;557 };558 void f(b g) { b e = g; }559}560 561// Implicit HD candidate competes with host candidate.562// a and b have implicit HD copy ctor. In copy ctor of b, ctor of a is resolved.563// copy ctor of a should win over a(short), otherwise there will be ambiguity564// due to conversion operator.565namespace TestImplicitHDWithH {566 struct a {567 a(short);568 __device__ operator unsigned() const;569 __device__ operator int() const;570 };571 struct b {572 a d;573 };574 void f(b g) { b e = g; }575}576 577// Implicit HD candidate competes with HD candidate.578// a and b have implicit HD copy ctor. In copy ctor of b, ctor of a is resolved.579// copy ctor of a should win over a(short), otherwise there will be ambiguity580// due to conversion operator.581namespace TestImplicitHDWithHD {582 struct a {583 __host__ __device__ a(short);584 __device__ operator unsigned() const;585 __device__ operator int() const;586 };587 struct b {588 a d;589 };590 void f(b g) { b e = g; }591}592 593// HD candidate competes with H candidate.594// HD has type mismatch whereas H has type match.595// In device compilation, H wins when -fgpu-defer-diag is off and HD wins596// when -fgpu-defer-diags is on. In both cases the diagnostic should be597// deferred.598namespace TestDeferNoMatchingFuncNotEmitted {599 template <typename> struct a {};600 namespace b {601 struct c : a<int> {};602 template <typename d> void ag(d);603 } // namespace b604 template <typename ae>605 __host__ __device__ void ag(a<ae>) {606 ae e;607 ag(e);608 }609 void f() { (void)ag<b::c>; }610}611 612namespace TestDeferNoMatchingFuncEmitted {613 template <typename> struct a {};614 namespace b {615 struct c : a<int> {};616 template <typename d> void ag(d);617 // devnodeferonly-note@-1{{'ag<TestDeferNoMatchingFuncEmitted::b::c>' declared here}}618 } // namespace b619 template <typename ae>620 __host__ __device__ void ag(a<ae>) {621 ae e;622 ag(e);623 // devnodeferonly-error@-1{{reference to __host__ function 'ag<TestDeferNoMatchingFuncEmitted::b::c>' in __host__ __device__ function}}624 // devdeferonly-error@-2{{no matching function for call to 'ag'}}625 // devdeferonly-note@-3{{called by 'ag<TestDeferNoMatchingFuncEmitted::b::c>'}}626 }627 __host__ __device__ void f() { (void)ag<b::c>; }628 // devnodeferonly-note@-1{{called by 'f'}}629 // devdeferonly-note@-2{{called by 'f'}}630}631 632// Two HD candidates compete with H candidate.633// HDs have type mismatch whereas H has type match.634// In device compilation, H wins when -fgpu-defer-diag is off and two HD win635// when -fgpu-defer-diags is on. In both cases the diagnostic should be636// deferred.637namespace TestDeferAmbiguityNotEmitted {638 template <typename> struct a {};639 namespace b {640 struct c : a<int> {};641 template <typename d> void ag(d, int);642 } // namespace b643 template <typename ae>644 __host__ __device__ void ag(a<ae>, float) {645 ae e;646 ag(e, 1);647 }648 template <typename ae>649 __host__ __device__ void ag(a<ae>, double) {650 }651 void f() {652 b::c x;653 ag(x, 1);654 }655}656 657namespace TestDeferAmbiguityEmitted {658 template <typename> struct a {};659 namespace b {660 struct c : a<int> {};661 template <typename d> void ag(d, int);662 // devnodeferonly-note@-1{{'ag<TestDeferAmbiguityEmitted::b::c>' declared here}}663 } // namespace b664 template <typename ae>665 __host__ __device__ void ag(a<ae>, float) {666 // devdeferonly-note@-1{{candidate function [with ae = int]}}667 ae e;668 ag(e, 1);669 }670 template <typename ae>671 __host__ __device__ void ag(a<ae>, double) {672 // devdeferonly-note@-1{{candidate function [with ae = int]}}673 }674 __host__ __device__ void f() {675 b::c x;676 ag(x, 1);677 // devnodeferonly-error@-1{{reference to __host__ function 'ag<TestDeferAmbiguityEmitted::b::c>' in __host__ __device__ function}}678 // devdeferonly-error@-2{{call to 'ag' is ambiguous}}679 }680}681 682// Implicit HD functions compute with H function and D function.683// In host compilation, foo(0.0, 2) should resolve to X::foo<double, int>.684// In device compilation, foo(0.0, 2) should resolve to foo(double, int).685// In either case there should be no ambiguity.686namespace TestImplicitHDWithHAndD {687 namespace X {688 inline double foo(double, double) { return 0;}689 inline constexpr float foo(float, float) { return 1;}690 inline constexpr long double foo(long double, long double) { return 2;}691 template<typename _Tp, typename _Up> inline constexpr double foo(_Tp, _Up) { return 3;}692 };693 using X::foo;694 inline __device__ double foo(double, double) { return 4;}695 inline __device__ float foo(float, int) { return 5;}696 inline __device__ float foo(int, int) { return 6;}697 inline __device__ double foo(double, int) { return 7;}698 inline __device__ float foo(float, float) { return 9;}699 template<typename _Tp, typename _Up> inline __device__ double foo(_Tp, _Up) { return 10;}700 701 int g() {702 return [](){703 return foo(0.0, 2);704 }();705 }706}707