test/SemaCUDA/function-overload.cu - third_party/swift-clang - Git at Google

 // 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 definition is here}}
 __host__ int hdh() { return 0; }            // expected-error {{redefinition of 'hdh'}}

 __host__ int hhd() { return 0; }            // expected-note {{previous definition is here}}
 __host__ __device__ int hhd() { return 0; } // expected-error {{redefinition of 'hhd'}}
 // expected-warning@-1 {{attribute declaration must precede definition}}
 // expected-note@-3 {{previous definition is here}}

 __host__ __device__ int hdd() { return 0; } // expected-note {{previous definition is here}}
 __device__ int hdd() { return 0; }          // expected-error {{redefinition of 'hdd'}}

 __device__ int dhd() { return 0; }          // expected-note {{previous definition is here}}
 __host__ __device__ int dhd() { return 0; } // expected-error {{redefinition of 'dhd'}}
 // expected-warning@-1 {{attribute declaration must precede definition}}
 // expected-note@-3 {{previous definition is here}}

 // Same tests for extern "C" functions.
 extern "C" __host__ int chh() { 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 definition is here}}
 extern "C" __host__ int chhd1() { return 0; }            // expected-error {{redefinition of 'chhd1'}}

 extern "C" __host__ int chhd2() { return 0; }            // expected-note {{previous definition is here}}
 extern "C" __host__ __device__ int chhd2() { return 0; } // expected-error {{redefinition of 'chhd2'}}
 // expected-warning@-1 {{attribute declaration must precede definition}}
 // expected-note@-3 {{previous definition is here}}

 // Helper functions to verify calling restrictions.
 __device__ 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();

   HostFnPtr fp_h = h;
   HostReturnTy ret_h = h();
   HostFnPtr fp_ch = ch;
   HostReturnTy ret_ch = ch();

   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();
   g<<<0,0>>>();
 #if !defined(__CUDA_ARCH__)
   // expected-error@-3 {{call to global function g not configured}}
 #else
   // expected-error@-5 {{no matching function for call to 'g'}}
   // expected-error@-5 {{reference to __global__ function 'g' in __host__ __device__ function}}
 #endif  // __CUDA_ARCH__
 }

 // 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 declaration is here}}
   __host__ ~d_h() {} // expected-error {{destructor cannot be redeclared}}
 };

 // H/D overloading is OK
 struct d_dh {
   __device__ ~d_dh() {}
   __host__ ~d_dh() {}
 };

 // HD is OK
 struct d_hd {
   __host__ __device__ ~d_hd() {}
 };

 // Mixing H/D and HD is not allowed.
 struct d_dhhd {
   __device__ ~d_dhhd() {}
   __host__ ~d_dhhd() {} // expected-note {{previous declaration is here}}
   __host__ __device__ ~d_dhhd() {} // expected-error {{destructor cannot be redeclared}}
 };

 struct d_hhd {
   __host__ ~d_hhd() {} // expected-note {{previous declaration is here}}
   __host__ __device__ ~d_hhd() {} // expected-error {{destructor cannot be redeclared}}
 };

 struct d_hdh {
   __host__ __device__ ~d_hdh() {} // expected-note {{previous declaration is here}}
   __host__ ~d_hdh() {} // expected-error {{destructor cannot be redeclared}}
 };

 struct d_dhd {
   __device__ ~d_dhd() {} // expected-note {{previous declaration is here}}
   __host__ __device__ ~d_dhd() {} // expected-error {{destructor cannot be redeclared}}
 };

 struct d_hdd {
   __host__ __device__ ~d_hdd() {} // expected-note {{previous declaration is here}}
   __device__ ~d_hdd() {} // expected-error {{destructor cannot be redeclared}}
 };

 // Test overloading of member functions
 struct m_h {
   void operator delete(void *ptr); // expected-note {{previous declaration is here}}
   __host__ void operator delete(void *ptr); // expected-error {{class member cannot be redeclared}}
 };

 // D/H overloading is OK
 struct m_dh {
   __device__ void operator delete(void *ptr);
   __host__ void operator delete(void *ptr);
 };

 // HD by itself is OK
 struct m_hd {
   __device__ __host__ void operator delete(void *ptr);
 };

 struct m_hhd {
   __host__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
   __host__ __device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
 };

 struct m_hdh {
   __host__ __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
   __host__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
 };

 struct m_dhd {
   __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
   __host__ __device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
 };

 struct m_hdd {
   __host__ __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
   __device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
 };

 // __global__ functions can't be overloaded based on attribute
 // difference.
 struct G {
   friend void friend_of_g(G &arg);
 private:
   int x;
 };
 __global__ void friend_of_g(G &arg) { int x = arg.x; } // expected-note {{previous definition is here}}
 void friend_of_g(G &arg) { int x = arg.x; } // expected-error {{redefinition of 'friend_of_g'}}

 // 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) {
   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);
 }

 __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(); }
 __host__ HostReturnTy host_only_function(int arg) { return HostReturnTy(); }
 __host__ HostReturnTy2 host_only_function(float arg) { return HostReturnTy2(); }

 __host__ __device__ void test_host_device_single_side_overloading() {
   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);
 }
	// 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 definition is here}}
	__host__ int hdh() { return 0; } // expected-error {{redefinition of 'hdh'}}

	__host__ int hhd() { return 0; } // expected-note {{previous definition is here}}
	__host__ __device__ int hhd() { return 0; } // expected-error {{redefinition of 'hhd'}}
	// expected-warning@-1 {{attribute declaration must precede definition}}
	// expected-note@-3 {{previous definition is here}}

	__host__ __device__ int hdd() { return 0; } // expected-note {{previous definition is here}}
	__device__ int hdd() { return 0; } // expected-error {{redefinition of 'hdd'}}

	__device__ int dhd() { return 0; } // expected-note {{previous definition is here}}
	__host__ __device__ int dhd() { return 0; } // expected-error {{redefinition of 'dhd'}}
	// expected-warning@-1 {{attribute declaration must precede definition}}
	// expected-note@-3 {{previous definition is here}}

	// Same tests for extern "C" functions.
	extern "C" __host__ int chh() { 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 definition is here}}
	extern "C" __host__ int chhd1() { return 0; } // expected-error {{redefinition of 'chhd1'}}

	extern "C" __host__ int chhd2() { return 0; } // expected-note {{previous definition is here}}
	extern "C" __host__ __device__ int chhd2() { return 0; } // expected-error {{redefinition of 'chhd2'}}
	// expected-warning@-1 {{attribute declaration must precede definition}}
	// expected-note@-3 {{previous definition is here}}

	// Helper functions to verify calling restrictions.
	__device__ 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();

	HostFnPtr fp_h = h;
	HostReturnTy ret_h = h();
	HostFnPtr fp_ch = ch;
	HostReturnTy ret_ch = ch();

	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();
	g<<<0,0>>>();
	#if !defined(__CUDA_ARCH__)
	// expected-error@-3 {{call to global function g not configured}}
	#else
	// expected-error@-5 {{no matching function for call to 'g'}}
	// expected-error@-5 {{reference to __global__ function 'g' in __host__ __device__ function}}
	#endif // __CUDA_ARCH__
	}

	// 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 declaration is here}}
	__host__ ~d_h() {} // expected-error {{destructor cannot be redeclared}}
	};

	// H/D overloading is OK
	struct d_dh {
	__device__ ~d_dh() {}
	__host__ ~d_dh() {}
	};

	// HD is OK
	struct d_hd {
	__host__ __device__ ~d_hd() {}
	};

	// Mixing H/D and HD is not allowed.
	struct d_dhhd {
	__device__ ~d_dhhd() {}
	__host__ ~d_dhhd() {} // expected-note {{previous declaration is here}}
	__host__ __device__ ~d_dhhd() {} // expected-error {{destructor cannot be redeclared}}
	};

	struct d_hhd {
	__host__ ~d_hhd() {} // expected-note {{previous declaration is here}}
	__host__ __device__ ~d_hhd() {} // expected-error {{destructor cannot be redeclared}}
	};

	struct d_hdh {
	__host__ __device__ ~d_hdh() {} // expected-note {{previous declaration is here}}
	__host__ ~d_hdh() {} // expected-error {{destructor cannot be redeclared}}
	};

	struct d_dhd {
	__device__ ~d_dhd() {} // expected-note {{previous declaration is here}}
	__host__ __device__ ~d_dhd() {} // expected-error {{destructor cannot be redeclared}}
	};

	struct d_hdd {
	__host__ __device__ ~d_hdd() {} // expected-note {{previous declaration is here}}
	__device__ ~d_hdd() {} // expected-error {{destructor cannot be redeclared}}
	};

	// Test overloading of member functions
	struct m_h {
	void operator delete(void *ptr); // expected-note {{previous declaration is here}}
	__host__ void operator delete(void *ptr); // expected-error {{class member cannot be redeclared}}
	};

	// D/H overloading is OK
	struct m_dh {
	__device__ void operator delete(void *ptr);
	__host__ void operator delete(void *ptr);
	};

	// HD by itself is OK
	struct m_hd {
	__device__ __host__ void operator delete(void *ptr);
	};

	struct m_hhd {
	__host__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
	__host__ __device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
	};

	struct m_hdh {
	__host__ __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
	__host__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
	};

	struct m_dhd {
	__device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
	__host__ __device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
	};

	struct m_hdd {
	__host__ __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
	__device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
	};

	// __global__ functions can't be overloaded based on attribute
	// difference.
	struct G {
	friend void friend_of_g(G &arg);
	private:
	int x;
	};
	__global__ void friend_of_g(G &arg) { int x = arg.x; } // expected-note {{previous definition is here}}
	void friend_of_g(G &arg) { int x = arg.x; } // expected-error {{redefinition of 'friend_of_g'}}

	// 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) {
	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);
	}

	__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(); }
	__host__ HostReturnTy host_only_function(int arg) { return HostReturnTy(); }
	__host__ HostReturnTy2 host_only_function(float arg) { return HostReturnTy2(); }

	__host__ __device__ void test_host_device_single_side_overloading() {
	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);
	}