Index: include/clang/Basic/LangOptions.def =================================================================== --- include/clang/Basic/LangOptions.def +++ include/clang/Basic/LangOptions.def @@ -175,6 +175,7 @@ LANGOPT(CUDADisableTargetCallChecks, 1, 0, "Disable checks for call targets (host, device, etc.)") LANGOPT(CUDATargetOverloads, 1, 0, "Enable function overloads based on CUDA target attributes") LANGOPT(CUDAAllowVariadicFunctions, 1, 0, "Allow variadic functions in CUDA device code") +LANGOPT(CUDARelaxedConstexpr, 1, 0, "Treat constexpr functions as __host__ __device__") LANGOPT(AssumeSaneOperatorNew , 1, 1, "implicit __attribute__((malloc)) for C++'s new operators") LANGOPT(SizedDeallocation , 1, 0, "enable sized deallocation functions") Index: include/clang/Driver/CC1Options.td =================================================================== --- include/clang/Driver/CC1Options.td +++ include/clang/Driver/CC1Options.td @@ -699,6 +699,8 @@ HelpText<"Enable function overloads based on CUDA target attributes.">; def fcuda_allow_variadic_functions : Flag<["-"], "fcuda-allow-variadic-functions">, HelpText<"Allow variadic functions in CUDA device code.">; +def fcuda_relaxed_constexpr : Flag<["-"], "fcuda-relaxed-constexpr">, + HelpText<"Treat constexpr functions as __host__ __device__.">; //===----------------------------------------------------------------------===// // OpenMP Options Index: lib/Driver/Tools.cpp =================================================================== --- lib/Driver/Tools.cpp +++ lib/Driver/Tools.cpp @@ -3594,6 +3594,7 @@ CmdArgs.push_back(Args.MakeArgString(AuxToolChain->getTriple().str())); CmdArgs.push_back("-fcuda-target-overloads"); CmdArgs.push_back("-fcuda-disable-target-call-checks"); + CmdArgs.push_back("-fcuda-relaxed-constexpr"); } if (Triple.isOSWindows() && (Triple.getArch() == llvm::Triple::arm || Index: lib/Frontend/CompilerInvocation.cpp =================================================================== --- lib/Frontend/CompilerInvocation.cpp +++ lib/Frontend/CompilerInvocation.cpp @@ -1569,6 +1569,9 @@ if (Args.hasArg(OPT_fcuda_allow_variadic_functions)) Opts.CUDAAllowVariadicFunctions = 1; + if (Args.hasArg(OPT_fcuda_relaxed_constexpr)) + Opts.CUDARelaxedConstexpr = 1; + if (Opts.ObjC1) { if (Arg *arg = Args.getLastArg(OPT_fobjc_runtime_EQ)) { StringRef value = arg->getValue(); Index: lib/Sema/SemaDecl.cpp =================================================================== --- lib/Sema/SemaDecl.cpp +++ lib/Sema/SemaDecl.cpp @@ -8006,6 +8006,15 @@ // Handle attributes. ProcessDeclAttributes(S, NewFD, D); + // With -fcuda-relaxed-constexpr, constexpr functions are treated as + // implicitly __host__ __device__. Device-side variadic functions are not + // allowed, so we just treat those as host-only. + if (getLangOpts().CUDA && NewFD->isConstexpr() && !NewFD->isVariadic() && + getLangOpts().CUDARelaxedConstexpr) { + NewFD->addAttr(CUDAHostAttr::CreateImplicit(Context)); + NewFD->addAttr(CUDADeviceAttr::CreateImplicit(Context)); + } + if (getLangOpts().OpenCL) { // OpenCL v1.1 s6.5: Using an address space qualifier in a function return // type declaration will generate a compilation error. Index: lib/Sema/SemaOverload.cpp =================================================================== --- lib/Sema/SemaOverload.cpp +++ lib/Sema/SemaOverload.cpp @@ -1126,13 +1126,10 @@ assert((OldTarget != CFT_InvalidTarget) && "Unexpected invalid target."); - // Don't allow mixing of HD with other kinds. This guarantees that - // we have only one viable function with this signature on any - // side of CUDA compilation . - // __global__ functions can't be overloaded based on attribute - // difference because, like HD, they also exist on both sides. - if ((NewTarget == CFT_HostDevice) || (OldTarget == CFT_HostDevice) || - (NewTarget == CFT_Global) || (OldTarget == CFT_Global)) + // Don't allow __global__ functions to be overloaded with other functions, + // based solely on their CUDA attributes. This guarantees that we have only + // one viable function with this signature on any side of CUDA compilation. + if ((NewTarget == CFT_Global) || (OldTarget == CFT_Global)) return false; // Allow overloading of functions with same signature, but Index: test/SemaCUDA/function-overload.cu =================================================================== --- test/SemaCUDA/function-overload.cu +++ test/SemaCUDA/function-overload.cu @@ -27,22 +27,18 @@ __host__ int dh(void) { return 2; } __device__ int dh(void) { return 2; } -// H/HD and D/HD are not allowed -__host__ __device__ int hdh(void) { return 5; } // expected-note {{previous definition is here}} -__host__ int hdh(void) { return 4; } // expected-error {{redefinition of 'hdh'}} +// H/HD and D/HD are OK +__host__ __device__ int hdh(void) { return 5; } +__host__ int hdh(void) { return 4; } -__host__ int hhd(void) { return 4; } // expected-note {{previous definition is here}} -__host__ __device__ int hhd(void) { return 5; } // expected-error {{redefinition of 'hhd'}} -// expected-warning@-1 {{attribute declaration must precede definition}} -// expected-note@-3 {{previous definition is here}} +__host__ int hhd(void) { return 4; } +__host__ __device__ int hhd(void) { return 5; } -__host__ __device__ int hdd(void) { return 7; } // expected-note {{previous definition is here}} -__device__ int hdd(void) { return 6; } // expected-error {{redefinition of 'hdd'}} +__host__ __device__ int hdd(void) { return 7; } +__device__ int hdd(void) { return 6; } -__device__ int dhd(void) { return 6; } // expected-note {{previous definition is here}} -__host__ __device__ int dhd(void) { return 7; } // expected-error {{redefinition of 'dhd'}} -// expected-warning@-1 {{attribute declaration must precede definition}} -// expected-note@-3 {{previous definition is here}} +__device__ int dhd(void) { return 6; } +__host__ __device__ int dhd(void) { return 7; } // Same tests for extern "C" functions extern "C" __host__ int chh(void) {return 11;} // expected-note {{previous definition is here}} @@ -52,14 +48,12 @@ extern "C" __device__ int cdh(void) {return 10;} extern "C" __host__ int cdh(void) {return 11;} -// H/HD and D/HD overloading is not allowed. -extern "C" __host__ __device__ int chhd1(void) {return 12;} // expected-note {{previous definition is here}} -extern "C" __host__ int chhd1(void) {return 13;} // expected-error {{redefinition of 'chhd1'}} +// H/HD and D/HD overloading is OK +extern "C" __host__ __device__ int chhd1(void) {return 12;} +extern "C" __host__ int chhd1(void) {return 13;} -extern "C" __host__ int chhd2(void) {return 13;} // expected-note {{previous definition is here}} -extern "C" __host__ __device__ int chhd2(void) {return 12;} // expected-error {{redefinition of 'chhd2'}} -// expected-warning@-1 {{attribute declaration must precede definition}} -// expected-note@-3 {{previous definition is here}} +extern "C" __host__ int chhd2(void) {return 13;} +extern "C" __host__ __device__ int chhd2(void) {return 12;} // Helper functions to verify calling restrictions. __device__ int d(void) { return 8; } @@ -71,10 +65,10 @@ __host__ void hostf(void) { fp_t dp = d; // expected-error@-1 {{reference to __device__ function 'd' in __host__ function}} - // expected-note@65 {{'d' declared here}} + // expected-note@59 {{'d' declared here}} fp_t cdp = cd; // expected-error@-1 {{reference to __device__ function 'cd' in __host__ function}} - // expected-note@68 {{'cd' declared here}} + // expected-note@62 {{'cd' declared here}} fp_t hp = h; fp_t chp = ch; fp_t dhp = dh; @@ -83,10 +77,10 @@ d(); // expected-error@-1 {{no matching function for call to 'd'}} - // expected-note@65 {{candidate function not viable: call to __device__ function from __host__ function}} + // expected-note@59 {{candidate function not viable: call to __device__ function from __host__ function}} cd(); // expected-error@-1 {{no matching function for call to 'cd'}} - // expected-note@68 {{candidate function not viable: call to __device__ function from __host__ function}} + // expected-note@62 {{candidate function not viable: call to __device__ function from __host__ function}} h(); ch(); dh(); @@ -101,27 +95,27 @@ fp_t cdp = cd; fp_t hp = h; // expected-error@-1 {{reference to __host__ function 'h' in __device__ function}} - // expected-note@66 {{'h' declared here}} + // expected-note@60 {{'h' declared here}} fp_t chp = ch; // expected-error@-1 {{reference to __host__ function 'ch' in __device__ function}} - // expected-note@69 {{'ch' declared here}} + // expected-note@63 {{'ch' declared here}} fp_t dhp = dh; fp_t cdhp = cdh; gp_t gp = g; // expected-error {{reference to __global__ function 'g' in __device__ function}} - // expected-note@67 {{'g' declared here}} + // expected-note@61 {{'g' declared here}} d(); cd(); h(); // expected-error {{no matching function for call to 'h'}} - // expected-note@66 {{candidate function not viable: call to __host__ function from __device__ function}} + // expected-note@60 {{candidate function not viable: call to __host__ function from __device__ function}} ch(); // expected-error {{no matching function for call to 'ch'}} - // expected-note@69 {{candidate function not viable: call to __host__ function from __device__ function}} + // expected-note@63 {{candidate function not viable: call to __host__ function from __device__ function}} dh(); cdh(); g(); // expected-error {{no matching function for call to 'g'}} - // expected-note@67 {{candidate function not viable: call to __global__ function from __device__ function}} + // expected-note@61 {{candidate function not viable: call to __global__ function from __device__ function}} g<<<0,0>>>(); // expected-error {{reference to __global__ function 'g' in __device__ function}} - // expected-note@67 {{'g' declared here}} + // expected-note@61 {{'g' declared here}} } __global__ void globalf(void) { @@ -129,30 +123,30 @@ fp_t cdp = cd; fp_t hp = h; // expected-error@-1 {{reference to __host__ function 'h' in __global__ function}} - // expected-note@66 {{'h' declared here}} + // expected-note@60 {{'h' declared here}} fp_t chp = ch; // expected-error@-1 {{reference to __host__ function 'ch' in __global__ function}} - // expected-note@69 {{'ch' declared here}} + // expected-note@63 {{'ch' declared here}} fp_t dhp = dh; fp_t cdhp = cdh; gp_t gp = g; // expected-error@-1 {{reference to __global__ function 'g' in __global__ function}} - // expected-note@67 {{'g' declared here}} + // expected-note@61 {{'g' declared here}} d(); cd(); h(); // expected-error@-1 {{no matching function for call to 'h'}} - // expected-note@66 {{candidate function not viable: call to __host__ function from __global__ function}} + // expected-note@60 {{candidate function not viable: call to __host__ function from __global__ function}} ch(); // expected-error@-1 {{no matching function for call to 'ch'}} - // expected-note@69 {{candidate function not viable: call to __host__ function from __global__ function}} + // expected-note@63 {{candidate function not viable: call to __host__ function from __global__ function}} dh(); cdh(); g(); // expected-error {{no matching function for call to 'g'}} - // expected-note@67 {{candidate function not viable: call to __global__ function from __global__ function}} + // expected-note@61 {{candidate function not viable: call to __global__ function from __global__ function}} g<<<0,0>>>(); // expected-error {{reference to __global__ function 'g' in __global__ function}} - // expected-note@67 {{'g' declared here}} + // expected-note@61 {{'g' declared here}} } __host__ __device__ void hostdevicef(void) { @@ -163,14 +157,14 @@ #if !defined(NOCHECKS) #if !defined(__CUDA_ARCH__) // expected-error@-6 {{reference to __device__ function 'd' in __host__ __device__ function}} - // expected-note@65 {{'d' declared here}} + // expected-note@59 {{'d' declared here}} // expected-error@-7 {{reference to __device__ function 'cd' in __host__ __device__ function}} - // expected-note@68 {{'cd' declared here}} + // expected-note@62 {{'cd' declared here}} #else // expected-error@-9 {{reference to __host__ function 'h' in __host__ __device__ function}} - // expected-note@66 {{'h' declared here}} + // expected-note@60 {{'h' declared here}} // expected-error@-10 {{reference to __host__ function 'ch' in __host__ __device__ function}} - // expected-note@69 {{'ch' declared here}} + // expected-note@63 {{'ch' declared here}} #endif #endif fp_t dhp = dh; @@ -178,7 +172,7 @@ gp_t gp = g; #if defined(__CUDA_ARCH__) // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}} - // expected-note@67 {{'g' declared here}} + // expected-note@61 {{'g' declared here}} #endif d(); @@ -188,14 +182,14 @@ #if !defined(NOCHECKS) #if !defined(__CUDA_ARCH__) // expected-error@-6 {{no matching function for call to 'd'}} - // expected-note@65 {{candidate function not viable: call to __device__ function from __host__ __device__ function}} + // expected-note@59 {{candidate function not viable: call to __device__ function from __host__ __device__ function}} // expected-error@-7 {{no matching function for call to 'cd'}} - // expected-note@68 {{candidate function not viable: call to __device__ function from __host__ __device__ function}} + // expected-note@62 {{candidate function not viable: call to __device__ function from __host__ __device__ function}} #else // expected-error@-9 {{no matching function for call to 'h'}} - // expected-note@66 {{candidate function not viable: call to __host__ function from __host__ __device__ function}} + // expected-note@60 {{candidate function not viable: call to __host__ function from __host__ __device__ function}} // expected-error@-10 {{no matching function for call to 'ch'}} - // expected-note@69 {{candidate function not viable: call to __host__ function from __host__ __device__ function}} + // expected-note@63 {{candidate function not viable: call to __host__ function from __host__ __device__ function}} #endif #endif @@ -207,9 +201,9 @@ // expected-error@-3 {{call to global function g not configured}} #else // expected-error@-5 {{no matching function for call to 'g'}} - // expected-note@67 {{candidate function not viable: call to __global__ function from __host__ __device__ function}} + // expected-note@61 {{candidate function not viable: call to __global__ function from __host__ __device__ function}} // expected-error@-6 {{reference to __global__ function 'g' in __host__ __device__ function}} - // expected-note@67 {{'g' declared here}} + // expected-note@61 {{'g' declared here}} #endif // __CUDA_ARCH__ } @@ -239,31 +233,31 @@ __host__ __device__ ~d_hd() {} }; -// Mixing H/D and HD is not allowed. +// Mixing H/D and HD is OK 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}} + __host__ ~d_dhhd() {} + __host__ __device__ ~d_dhhd() {} }; struct d_hhd { - __host__ ~d_hhd() {} // expected-note {{previous declaration is here}} - __host__ __device__ ~d_hhd() {} // expected-error {{destructor cannot be redeclared}} + __host__ ~d_hhd() {} + __host__ __device__ ~d_hhd() {} }; struct d_hdh { - __host__ __device__ ~d_hdh() {} // expected-note {{previous declaration is here}} - __host__ ~d_hdh() {} // expected-error {{destructor cannot be redeclared}} + __host__ __device__ ~d_hdh() {} + __host__ ~d_hdh() {} }; struct d_dhd { - __device__ ~d_dhd() {} // expected-note {{previous declaration is here}} - __host__ __device__ ~d_dhd() {} // expected-error {{destructor cannot be redeclared}} + __device__ ~d_dhd() {} + __host__ __device__ ~d_dhd() {} }; struct d_hdd { - __host__ __device__ ~d_hdd() {} // expected-note {{previous declaration is here}} - __device__ ~d_hdd() {} // expected-error {{destructor cannot be redeclared}} + __host__ __device__ ~d_hdd() {} + __device__ ~d_hdd() {} }; // Test overloading of member functions @@ -284,23 +278,23 @@ }; 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}} + __host__ void operator delete(void *ptr) {} + __host__ __device__ void operator delete(void *ptr) {} }; 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}} + __host__ __device__ void operator delete(void *ptr) {} + __host__ void operator delete(void *ptr) {} }; 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}} + __device__ void operator delete(void *ptr) {} + __host__ __device__ void operator delete(void *ptr) {} }; 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}} + __host__ __device__ void operator delete(void *ptr) {} + __device__ void operator delete(void *ptr) {} }; // __global__ functions can't be overloaded based on attribute Index: test/SemaCUDA/relaxed-constexpr.cu =================================================================== --- /dev/null +++ test/SemaCUDA/relaxed-constexpr.cu @@ -0,0 +1,23 @@ +// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify %s -fcuda-target-overloads -fcuda-relaxed-constexpr +// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify %s -fcuda-target-overloads -fcuda-relaxed-constexpr -fcuda-is-device + +// expected-no-diagnostics + +#include "Inputs/cuda.h" + +static __device__ void f1(); +constexpr void f1(); + +__device__ void f2(); +static constexpr void f2(); + +// Different potential error depending on the order of declaration. +constexpr void f3(); +static __device__ void f3(); + +static constexpr void f4(); +__device__ void f4(); + +// Variadic device functions are not allowed, so this is just treated as +// host-only. +constexpr void variadic(const char*, ...);