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*, ...);