[CUDA] Merge most of CodeGenCUDA/function-overload.cu into SemaCUDA/function-overload.cu.

[jlebar]

Summary

Previously we were using the codegen test to ensure that we choose the
right overload. But we can do this within sema, with a bit of
cleverness.

I left the constructor/destructor checks in CodeGen, because these
overloads (particularly on the destructors) are hard to check in Sema.

test/CodeGenCUDA/function-overload.cu

 // REQUIRES: x86-registered-target
 // REQUIRES: nvptx-registered-target
 
-// Make sure we handle target overloads correctly.
+// Make sure we handle target overloads correctly.  Most of this is checked in
+// sema, but special functions like constructors and destructors are here.
+//
 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu \
 // RUN:     -fcuda-target-overloads -emit-llvm -o - %s \
 // RUN:     | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-HOST %s
 // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device \
 // RUN:     -fcuda-target-overloads -emit-llvm -o - %s \
 // RUN:     | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-DEVICE \
 // RUN:       -check-prefix=CHECK-DEVICE-STRICT %s
 
 // Check target overloads handling with disabled call target checks.
 // RUN: %clang_cc1 -DNOCHECKS -triple x86_64-unknown-linux-gnu -emit-llvm \
 // RUN:    -fcuda-disable-target-call-checks -fcuda-target-overloads -o - %s \
 // RUN:     | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-HOST \
 // RUN:    -check-prefix=CHECK-BOTH-NC -check-prefix=CHECK-HOST-NC %s
 // RUN: %clang_cc1 -DNOCHECKS -triple nvptx64-nvidia-cuda -emit-llvm \
 // RUN:    -fcuda-disable-target-call-checks -fcuda-target-overloads \
 // RUN:    -fcuda-is-device -o - %s \
 // RUN:     | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-DEVICE \
 // RUN:    -check-prefix=CHECK-BOTH-NC -check-prefix=CHECK-DEVICE-NC %s
 
 #include "Inputs/cuda.h"
 
-typedef int (*fp_t)(void);
-typedef void (*gp_t)(void);
-
-// CHECK-HOST: @hp = global i32 ()* @_Z1hv
-// CHECK-HOST: @chp = global i32 ()* @ch
-// CHECK-HOST: @dhp = global i32 ()* @_Z2dhv
-// CHECK-HOST: @cdhp = global i32 ()* @cdh
-// CHECK-HOST: @gp = global void ()* @_Z1gv
-
-// CHECK-BOTH-LABEL: define i32 @_Z2dhv()
-__device__ int dh(void) { return 1; }
-// CHECK-DEVICE: ret i32 1
-__host__ int dh(void) { return 2; }
-// CHECK-HOST:   ret i32 2
-
-// CHECK-BOTH-LABEL: define i32 @_Z2hdv()
-__host__ __device__ int hd(void) { return 3; }
-// CHECK-BOTH:   ret i32 3
-
-// CHECK-DEVICE-LABEL: define i32 @_Z1dv()
-__device__ int d(void) { return 8; }
-// CHECK-DEVICE:   ret i32 8
-
-// CHECK-HOST-LABEL: define i32 @_Z1hv()
-__host__ int h(void) { return 9; }
-// CHECK-HOST:   ret i32 9
-
-// CHECK-BOTH-LABEL: define void @_Z1gv()
-__global__ void g(void) {}
-// CHECK-BOTH:   ret void
-
-// mangled names of extern "C" __host__ __device__ functions clash
-// with those of their __host__/__device__ counterparts, so
-// overloading of extern "C" functions can only happen for __host__
-// and __device__ functions -- we never codegen them in the same
-// compilation and therefore mangled name conflict is not a problem.
-
-// CHECK-BOTH-LABEL: define i32 @cdh()
-extern "C" __device__ int cdh(void) {return 10;}
-// CHECK-DEVICE:   ret i32 10
-extern "C" __host__ int cdh(void) {return 11;}
-// CHECK-HOST:     ret i32 11
-
-// CHECK-DEVICE-LABEL: define i32 @cd()
-extern "C" __device__ int cd(void) {return 12;}
-// CHECK-DEVICE:   ret i32 12
-
-// CHECK-HOST-LABEL: define i32 @ch()
-extern "C" __host__ int ch(void) {return 13;}
-// CHECK-HOST:     ret i32 13
-
-// CHECK-BOTH-LABEL: define i32 @chd()
-extern "C" __host__ __device__ int chd(void) {return 14;}
-// CHECK-BOTH:     ret i32 14
-
-// 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> T template_vs_function(T arg) { return 15; }
-__device__ float template_vs_function(float arg) { return 16; }
-
-// 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.
-// CHECK-BOTH-LABEL: define void @_Z5hd_tfv()
-__host__ __device__ void hd_tf(void) {
-  template_vs_function(1.0f);
-  // CHECK-HOST: call float @_Z20template_vs_functionIfET_S0_(float
-  // CHECK-DEVICE: call float @_Z20template_vs_functionf(float
-  template_vs_function(2.0);
-  // CHECK-HOST: call double @_Z20template_vs_functionIdET_S0_(double
-  // CHECK-DEVICE: call float @_Z20template_vs_functionf(float
-}
-
-// Calls from __host__ and __device__ functions should always call the
-// overloaded function that matches their mode.
-// CHECK-HOST-LABEL: define void @_Z4h_tfv()
-__host__ void h_tf() {
-  template_vs_function(1.0f);
-  // CHECK-HOST: call float @_Z20template_vs_functionIfET_S0_(float
-  template_vs_function(2.0);
-  // CHECK-HOST: call double @_Z20template_vs_functionIdET_S0_(double
-}
-
-// CHECK-DEVICE-LABEL: define void @_Z4d_tfv()
-__device__ void d_tf() {
-  template_vs_function(1.0f);
-  // CHECK-DEVICE: call float @_Z20template_vs_functionf(float
-  template_vs_function(2.0);
-  // CHECK-DEVICE: call float @_Z20template_vs_functionf(float
-}
-
-// In case 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> T template_vs_hd_function(T arg) { return 15; }
-__host__ __device__ float template_vs_hd_function(float arg) { return 16; }
-
-// CHECK-BOTH-LABEL: define void @_Z7hd_thdfv()
-__host__ __device__ void hd_thdf() {
-  template_vs_hd_function(1.0f);
-  // CHECK-HOST: call float @_Z23template_vs_hd_functionf(float
-  // CHECK-DEVICE: call float @_Z23template_vs_hd_functionf(float
-  template_vs_hd_function(1);
-  // CHECK-HOST: call i32 @_Z23template_vs_hd_functionIiET_S0_(i32
-  // CHECK-DEVICE-STRICT: call float @_Z23template_vs_hd_functionf(float
-  // CHECK-DEVICE-NC: call i32 @_Z23template_vs_hd_functionIiET_S0_(i32
-}
-
-// CHECK-HOST-LABEL: define void @_Z6h_thdfv()
-__host__ void h_thdf() {
-  template_vs_hd_function(1.0f);
-  // CHECK-HOST: call float @_Z23template_vs_hd_functionf(float
-  template_vs_hd_function(1);
-  // CHECK-HOST: call i32 @_Z23template_vs_hd_functionIiET_S0_(i32
-}
-
-// CHECK-DEVICE-LABEL: define void @_Z6d_thdfv()
-__device__ void d_thdf() {
-  template_vs_hd_function(1.0f);
-  // CHECK-DEVICE: call float @_Z23template_vs_hd_functionf(float
-  template_vs_hd_function(1);
-  // Host-only function template is not callable with strict call checks,
-  // so for device side HD function will be the only choice.
-  // CHECK-DEVICE: call float @_Z23template_vs_hd_functionf(float
-}
-
-// 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__ float device_only_function(int arg) { return 17; }
-__device__ float device_only_function(float arg) { return 18; }
-
-__host__ float host_only_function(int arg) { return 19; }
-__host__ float host_only_function(float arg) { return 20; }
-
-// CHECK-BOTH-LABEL: define void @_Z6hd_dofv()
-__host__ __device__ void hd_dof() {
-#ifdef NOCHECKS
-  device_only_function(1.0f);
-  // CHECK-BOTH-NC: call float @_Z20device_only_functionf(float
-  device_only_function(1);
-  // CHECK-BOTH-NC: call float @_Z20device_only_functioni(i32
-  host_only_function(1.0f);
-  // CHECK-BOTH-NC: call float @_Z18host_only_functionf(float
-  host_only_function(1);
-  // CHECK-BOTH-NC: call float @_Z18host_only_functioni(i32
-#endif
-}
-
-
-// CHECK-HOST-LABEL: define void @_Z5hostfv()
-__host__ void hostf(void) {
-  fp_t hp = h; // CHECK-HOST: store {{.*}} @_Z1hv, {{.*}} %hp,
-  fp_t chp = ch; // CHECK-HOST: store {{.*}} @ch, {{.*}} %chp,
-  fp_t dhp = dh; // CHECK-HOST: store {{.*}} @_Z2dhv, {{.*}} %dhp,
-  fp_t cdhp = cdh; // CHECK-HOST: store {{.*}} @cdh, {{.*}} %cdhp,
-  fp_t hdp = hd; // CHECK-HOST: store {{.*}} @_Z2hdv, {{.*}} %hdp,
-  fp_t chdp = chd; // CHECK-HOST: store {{.*}} @chd, {{.*}} %chdp,
-  gp_t gp = g; // CHECK-HOST: store {{.*}} @_Z1gv, {{.*}} %gp,
-
-  h();     // CHECK-HOST: call i32 @_Z1hv()
-  ch();    // CHECK-HOST: call i32 @ch()
-  dh();    // CHECK-HOST: call i32 @_Z2dhv()
-  cdh();   // CHECK-HOST: call i32 @cdh()
-  g<<<0,0>>>();  // CHECK-HOST: call void @_Z1gv()
-}
-
-// CHECK-DEVICE-LABEL: define void @_Z7devicefv()
-__device__ void devicef(void) {
-  fp_t dp = d;   // CHECK-DEVICE: store {{.*}} @_Z1dv, {{.*}} %dp,
-  fp_t cdp = cd; // CHECK-DEVICE: store {{.*}} @cd, {{.*}} %cdp,
-  fp_t dhp = dh; // CHECK-DEVICE: store {{.*}} @_Z2dhv, {{.*}} %dhp,
-  fp_t cdhp = cdh; // CHECK-DEVICE: store {{.*}} @cdh, {{.*}} %cdhp,
-  fp_t hdp = hd; // CHECK-DEVICE: store {{.*}} @_Z2hdv, {{.*}} %hdp,
-  fp_t chdp = chd; // CHECK-DEVICE: store {{.*}} @chd, {{.*}} %chdp,
-
-  d();     // CHECK-DEVICE: call i32 @_Z1dv()
-  cd();    // CHECK-DEVICE: call i32 @cd()
-  dh();    // CHECK-DEVICE: call i32 @_Z2dhv()
-  cdh();   // CHECK-DEVICE: call i32 @cdh()
-}
-
-// CHECK-BOTH-LABEL: define void @_Z11hostdevicefv()
-__host__ __device__ void hostdevicef(void) {
-#if defined (NOCHECKS)
-  fp_t dp = d;   // CHECK-BOTH-NC: store {{.*}} @_Z1dv, {{.*}} %dp,
-  fp_t cdp = cd; // CHECK-BOTH-NC: store {{.*}} @cd, {{.*}} %cdp,
-  fp_t hp = h; // CHECK-BOTH-NC: store {{.*}} @_Z1hv, {{.*}} %hp,
-  fp_t chp = ch; // CHECK-BOTH-NC: store {{.*}} @ch, {{.*}} %chp,
-#endif
-  fp_t dhp = dh; // CHECK-BOTH: store {{.*}} @_Z2dhv, {{.*}} %dhp,
-  fp_t cdhp = cdh; // CHECK-BOTH: store {{.*}} @cdh, {{.*}} %cdhp,
-  fp_t hdp = hd; // CHECK-BOTH: store {{.*}} @_Z2hdv, {{.*}} %hdp,
-  fp_t chdp = chd; // CHECK-BOTH: store {{.*}} @chd, {{.*}} %chdp,
-#if defined (NOCHECKS) && !defined(__CUDA_ARCH__)
-  gp_t gp = g; // CHECK-HOST-NC: store {{.*}} @_Z1gv, {{.*}} %gp,
-#endif
-
-#if defined (NOCHECKS)
-  d();     // CHECK-BOTH-NC: call i32 @_Z1dv()
-  cd();    // CHECK-BOTH-NC: call i32 @cd()
-  h();     // CHECK-BOTH-NC: call i32 @_Z1hv()
-  ch();    // CHECK-BOTH-NC: call i32 @ch()
-#endif
-  dh();    // CHECK-BOTH: call i32 @_Z2dhv()
-  cdh();   // CHECK-BOTH: call i32 @cdh()
-#if defined (NOCHECKS) && !defined(__CUDA_ARCH__)
-  g<<<0,0>>>();  // CHECK-HOST-NC: call void @_Z1gv()
-#endif
-}
-
-// Test for address of overloaded function resolution in the global context.
-fp_t hp = h;
-fp_t chp = ch;
-fp_t dhp = dh;
-fp_t cdhp = cdh;
-gp_t gp = g;
-
-int x;
 // Check constructors/destructors for D/H functions
+int x;
 struct s_cd_dh {
   __host__ s_cd_dh() { x = 11; }
   __device__ s_cd_dh() { x = 12; }
   __host__ ~s_cd_dh() { x = 21; }
   __device__ ~s_cd_dh() { x = 22; }
 };
 
 struct s_cd_hd {

 // CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_hdD2Ev(
 // CHECK-BOTH: store i32 32,
 // CHECK-BOTH: ret void
 
 // CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_dhD2Ev(
 // CHECK-HOST:   store i32 21,
 // CHECK-DEVICE: store i32 22,
 // CHECK-BOTH: ret void
-

test/SemaCUDA/function-overload.cu

 // RUN: %clang_cc1 -DNOCHECKS -triple x86_64-unknown-linux-gnu -fsyntax-only \
 // RUN:    -fcuda-disable-target-call-checks -fcuda-target-overloads -verify %s
 // RUN: %clang_cc1 -DNOCHECKS -triple nvptx64-nvidia-cuda -fsyntax-only \
 // RUN:    -fcuda-disable-target-call-checks -fcuda-target-overloads \
 // RUN:    -fcuda-is-device -verify %s
 
 #include "Inputs/cuda.h"
 
-typedef int (*fp_t)();
-typedef void (*gp_t)();
+// 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__ int dh() { return 2; }
-__device__ int dh() { return 2; }
+__host__ HostReturnTy dh() { return HostReturnTy(); }
+__device__ DeviceReturnTy dh() { return DeviceReturnTy(); }
 
 // H/HD and D/HD are not allowed.
-__host__ __device__ int hdh() { return 5; } // expected-note {{previous definition is here}}
-__host__ int hdh() { return 4; } // expected-error {{redefinition of 'hdh'}}
+__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 4; } // expected-note {{previous definition is here}}
-__host__ __device__ int hhd() { return 5; } // expected-error {{redefinition of 'hhd'}}
+__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 7; } // expected-note {{previous definition is here}}
-__device__ int hdd() { return 6; } // expected-error {{redefinition of 'hdd'}}
+__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 6; } // expected-note {{previous definition is here}}
-__host__ __device__ int dhd() { return 7; } // expected-error {{redefinition of 'dhd'}}
+__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 11;} // expected-note {{previous definition is here}}
-extern "C" int chh() {return 11;} // expected-error {{redefinition of 'chh'}}
+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__ int cdh() {return 10;}
-extern "C" __host__ int cdh() {return 11;}
+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 12;} // expected-note {{previous definition is here}}
-extern "C" __host__ int chhd1() {return 13;} // expected-error {{redefinition of 'chhd1'}}
+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 13;} // expected-note {{previous definition is here}}
-extern "C" __host__ __device__ int chhd2() {return 12;} // expected-error {{redefinition of 'chhd2'}}
+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__ int d() { return 8; }
+__device__ DeviceReturnTy d() { return DeviceReturnTy(); }
 // expected-note@-1 0+ {{'d' declared here}}
 // expected-note@-2 0+ {{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__ int h() { return 9; }
+__host__ HostReturnTy h() { return HostReturnTy(); }
 // expected-note@-1 0+ {{'h' declared here}}
 // expected-note@-2 0+ {{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 0+ {{candidate function not viable: call to __host__ function from __global__ function}}
 
 __global__ void g() {}
 // expected-note@-1 0+ {{'g' declared here}}
 // expected-note@-2 0+ {{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 0+ {{candidate function not viable: call to __global__ function from __global__ function}}
 
-extern "C" __device__ int cd() {return 10;}
+extern "C" __device__ DeviceReturnTy cd() { return DeviceReturnTy(); }
 // expected-note@-1 0+ {{'cd' declared here}}
 // expected-note@-2 0+ {{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__ int ch() {return 11;}
+extern "C" __host__ HostReturnTy ch() { return HostReturnTy(); }
 // expected-note@-1 0+ {{'ch' declared here}}
 // expected-note@-2 0+ {{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 0+ {{candidate function not viable: call to __host__ function from __global__ function}}
 
 __host__ void hostf() {
-  fp_t dp = d;
-  // expected-error@-1 {{reference to __device__ function 'd' in __host__ function}}
-  fp_t cdp = cd;
-  // expected-error@-1 {{reference to __device__ function 'cd' in __host__ function}}
-  fp_t hp = h;
-  fp_t chp = ch;
-  fp_t dhp = dh;
-  fp_t cdhp = cdh;
-  gp_t gp = g;
-
-  d();
-  // expected-error@-1 {{no matching function for call to 'd'}}
-  cd();
-  // expected-error@-1 {{no matching function for call to 'cd'}}
-  h();
-  ch();
-  dh();
-  cdh();
+  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>>>();
+  g<<<0, 0>>>();
 }
 
 __device__ void devicef() {
-  fp_t dp = d;
-  fp_t cdp = cd;
-  fp_t hp = h;
-  // expected-error@-1 {{reference to __host__ function 'h' in __device__ function}}
-  fp_t chp = ch;
-  // expected-error@-1 {{reference to __host__ function 'ch' in __device__ function}}
-  fp_t dhp = dh;
-  fp_t cdhp = cdh;
-  gp_t gp = g; // expected-error {{reference to __global__ function 'g' in __device__ function}}
-
-  d();
-  cd();
-  h(); // expected-error {{no matching function for call to 'h'}}
-  ch(); // expected-error {{no matching function for call to 'ch'}}
-  dh();
-  cdh();
+  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() {
-  fp_t dp = d;
-  fp_t cdp = cd;
-  fp_t hp = h;
-  // expected-error@-1 {{reference to __host__ function 'h' in __global__ function}}
-  fp_t chp = ch;
-  // expected-error@-1 {{reference to __host__ function 'ch' in __global__ function}}
-  fp_t dhp = dh;
-  fp_t cdhp = cdh;
-  gp_t gp = g;
-  // expected-error@-1 {{reference to __global__ function 'g' in __global__ function}}
-
-  d();
-  cd();
-  h();
-  // expected-error@-1 {{no matching function for call to 'h'}}
-  ch();
-  // expected-error@-1 {{no matching function for call to 'ch'}}
-  dh();
-  cdh();
+  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() {
-  fp_t dp = d;
-  fp_t cdp = cd;
+  DeviceFnPtr fp_d = d;
+  DeviceReturnTy ret_d = d();
+  DeviceFnPtr fp_cd = cd;
+  DeviceReturnTy ret_cd = cd();
 #if !defined(NOCHECKS) && !defined(__CUDA_ARCH__)
-  // expected-error@-3 {{reference to __device__ function 'd' in __host__ __device__ function}}
-  // expected-error@-3 {{reference to __device__ function 'cd' in __host__ __device__ function}}
+  // expected-error@-5 {{reference to __device__ function 'd' in __host__ __device__ function}}
+  // expected-error@-5 {{no matching function for call to 'd'}}
+  // expected-error@-5 {{reference to __device__ function 'cd' in __host__ __device__ function}}
+  // expected-error@-5 {{no matching function for call to 'cd'}}
 #endif
 
-  fp_t hp = h;
-  fp_t chp = ch;
+  HostFnPtr fp_h = h;
+  HostReturnTy ret_h = h();
+  HostFnPtr fp_ch = ch;
+  HostReturnTy ret_ch = ch();
 #if !defined(NOCHECKS) && defined(__CUDA_ARCH__)
-  // expected-error@-3 {{reference to __host__ function 'h' in __host__ __device__ function}}
-  // expected-error@-3 {{reference to __host__ function 'ch' in __host__ __device__ function}}
+  // expected-error@-5 {{reference to __host__ function 'h' in __host__ __device__ function}}
+  // expected-error@-5 {{no matching function for call to 'h'}}
+  // expected-error@-5 {{reference to __host__ function 'ch' in __host__ __device__ function}}
+  // expected-error@-5 {{no matching function for call to 'ch'}}
 #endif
 
-  fp_t dhp = dh;
-  fp_t cdhp = cdh;
-  gp_t gp = g;
+  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
-
-  d();
-  cd();
-#if !defined(NOCHECKS) && !defined(__CUDA_ARCH__)
-  // expected-error@-3 {{no matching function for call to 'd'}}
-  // expected-error@-3 {{no matching function for call to 'cd'}}
-#endif
-
-  h();
-  ch();
-#if !defined(NOCHECKS) && defined(__CUDA_ARCH__)
-  // expected-error@-3 {{no matching function for call to 'h'}}
-  // expected-error@-3 {{no matching function for call to 'ch'}}
-#endif
-
-  dh();
-  cdh();
   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.
-fp_t hp = h;
-fp_t chp = ch;
-fp_t dhp = dh;
-fp_t cdhp = cdh;
-gp_t gp = g;
+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}}
 };

 // 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);
+
+#if defined(__CUDA_ARCH__) && !defined(NOCHECKS)
+  typedef HostDeviceReturnTy ExpectedReturnTy;
+#else
+  typedef TemplateReturnTy ExpectedReturnTy;
+#endif
+  ExpectedReturnTy 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_nochecks_overloading() {
+#ifdef NOCHECKS
+  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);
+#endif
+}