r264207 - [CUDA] Merge most of CodeGenCUDA/function-overload.cu into SemaCUDA/function-overload.cu.
Justin Lebar via cfe-commits
cfe-commits at lists.llvm.org
Wed Mar 23 15:42:30 PDT 2016
Author: jlebar
Date: Wed Mar 23 17:42:30 2016
New Revision: 264207
URL: http://llvm.org/viewvc/llvm-project?rev=264207&view=rev
Log:
[CUDA] Merge most of CodeGenCUDA/function-overload.cu into SemaCUDA/function-overload.cu.
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.
Reviewers: tra
Subscribers: cfe-commits
Differential Revision: http://reviews.llvm.org/D18386
Modified:
cfe/trunk/test/CodeGenCUDA/function-overload.cu
cfe/trunk/test/SemaCUDA/function-overload.cu
Modified: cfe/trunk/test/CodeGenCUDA/function-overload.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/function-overload.cu?rev=264207&r1=264206&r2=264207&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenCUDA/function-overload.cu (original)
+++ cfe/trunk/test/CodeGenCUDA/function-overload.cu Wed Mar 23 17:42:30 2016
@@ -1,7 +1,9 @@
// 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
@@ -23,235 +25,8 @@
#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; }
@@ -300,4 +75,3 @@ void wrapper() {
// CHECK-HOST: store i32 21,
// CHECK-DEVICE: store i32 22,
// CHECK-BOTH: ret void
-
Modified: cfe/trunk/test/SemaCUDA/function-overload.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/function-overload.cu?rev=264207&r1=264206&r2=264207&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/function-overload.cu (original)
+++ cfe/trunk/test/SemaCUDA/function-overload.cu Wed Mar 23 17:42:30 2016
@@ -16,58 +16,80 @@
#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 at -1 {{attribute declaration must precede definition}}
// expected-note at -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 at -1 {{attribute declaration must precede definition}}
// expected-note at -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 at -1 {{attribute declaration must precede definition}}
// expected-note at -3 {{previous definition is here}}
// Helper functions to verify calling restrictions.
-__device__ int d() { return 8; }
+__device__ DeviceReturnTy d() { return DeviceReturnTy(); }
// expected-note at -1 1+ {{'d' declared here}}
// expected-note at -2 1+ {{candidate function not viable: call to __device__ function from __host__ function}}
// expected-note at -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 at -1 1+ {{'h' declared here}}
// expected-note at -2 1+ {{candidate function not viable: call to __host__ function from __device__ function}}
// expected-note at -3 0+ {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
@@ -79,123 +101,112 @@ __global__ void g() {}
// expected-note at -3 0+ {{candidate function not viable: call to __global__ function from __host__ __device__ function}}
// expected-note at -4 1+ {{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 at -1 1+ {{'cd' declared here}}
// expected-note at -2 1+ {{candidate function not viable: call to __device__ function from __host__ function}}
// expected-note at -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 at -1 1+ {{'ch' declared here}}
// expected-note at -2 1+ {{candidate function not viable: call to __host__ function from __device__ function}}
// expected-note at -3 0+ {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
// expected-note at -4 1+ {{candidate function not viable: call to __host__ function from __global__ function}}
__host__ void hostf() {
- fp_t dp = d;
- // expected-error at -1 {{reference to __device__ function 'd' in __host__ function}}
- fp_t cdp = cd;
- // expected-error at -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 at -1 {{no matching function for call to 'd'}}
- cd();
- // expected-error at -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 at -1 {{reference to __host__ function 'h' in __device__ function}}
- fp_t chp = ch;
- // expected-error at -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 at -1 {{reference to __host__ function 'h' in __global__ function}}
- fp_t chp = ch;
- // expected-error at -1 {{reference to __host__ function 'ch' in __global__ function}}
- fp_t dhp = dh;
- fp_t cdhp = cdh;
- gp_t gp = g;
- // expected-error at -1 {{reference to __global__ function 'g' in __global__ function}}
-
- d();
- cd();
- h();
- // expected-error at -1 {{no matching function for call to 'h'}}
- ch();
- // expected-error at -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 at -3 {{reference to __device__ function 'd' in __host__ __device__ function}}
- // expected-error at -3 {{reference to __device__ function 'cd' in __host__ __device__ function}}
+ // expected-error at -5 {{reference to __device__ function 'd' in __host__ __device__ function}}
+ // expected-error at -5 {{no matching function for call to 'd'}}
+ // expected-error at -5 {{reference to __device__ function 'cd' in __host__ __device__ function}}
+ // expected-error at -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 at -3 {{reference to __host__ function 'h' in __host__ __device__ function}}
- // expected-error at -3 {{reference to __host__ function 'ch' in __host__ __device__ function}}
+ // expected-error at -5 {{reference to __host__ function 'h' in __host__ __device__ function}}
+ // expected-error at -5 {{no matching function for call to 'h'}}
+ // expected-error at -5 {{reference to __host__ function 'ch' in __host__ __device__ function}}
+ // expected-error at -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 at -2 {{reference to __global__ function 'g' in __host__ __device__ function}}
#endif
-
- d();
- cd();
-#if !defined(NOCHECKS) && !defined(__CUDA_ARCH__)
- // expected-error at -3 {{no matching function for call to 'd'}}
- // expected-error at -3 {{no matching function for call to 'cd'}}
-#endif
-
- h();
- ch();
-#if !defined(NOCHECKS) && defined(__CUDA_ARCH__)
- // expected-error at -3 {{no matching function for call to 'h'}}
- // expected-error at -3 {{no matching function for call to 'ch'}}
-#endif
-
- dh();
- cdh();
g();
g<<<0,0>>>();
#if !defined(__CUDA_ARCH__)
@@ -207,11 +218,11 @@ __host__ __device__ void hostdevicef() {
}
// 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
@@ -305,3 +316,96 @@ private:
};
__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
+}
More information about the cfe-commits
mailing list