Index: include/clang/Basic/LangOptions.def
===================================================================
--- include/clang/Basic/LangOptions.def
+++ include/clang/Basic/LangOptions.def
@@ -174,6 +174,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, "Allow calls to constexpr functions from CUDA device code")
 
 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
@@ -697,6 +697,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<"Allow calls to constexpr functions from CUDA device code.">;
 
 //===----------------------------------------------------------------------===//
 // OpenMP Options
Index: include/clang/Driver/Options.td
===================================================================
--- include/clang/Driver/Options.td
+++ include/clang/Driver/Options.td
@@ -382,6 +382,8 @@
   HelpText<"Enable device-side debug info generation. Disables ptxas optimizations.">;
 def cuda_path_EQ : Joined<["--"], "cuda-path=">, Group<i_Group>,
   HelpText<"CUDA installation path">;
+def cuda_relaxed_constexpr : Flag<["--"], "cuda-relaxed-constexpr">,
+  HelpText<"Allow calls to constexpr functions from CUDA device code">;
 def dA : Flag<["-"], "dA">, Group<d_Group>;
 def dD : Flag<["-"], "dD">, Group<d_Group>, Flags<[CC1Option]>,
   HelpText<"Print macro definitions in -E mode in addition to normal output">;
Index: lib/Driver/Tools.cpp
===================================================================
--- lib/Driver/Tools.cpp
+++ lib/Driver/Tools.cpp
@@ -3594,6 +3594,9 @@
     CmdArgs.push_back(Args.MakeArgString(AuxToolChain->getTriple().str()));
     CmdArgs.push_back("-fcuda-target-overloads");
     CmdArgs.push_back("-fcuda-disable-target-call-checks");
+
+    if (Args.hasArg(options::OPT_cuda_relaxed_constexpr))
+      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
@@ -1576,6 +1576,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/SemaCUDA.cpp
===================================================================
--- lib/Sema/SemaCUDA.cpp
+++ lib/Sema/SemaCUDA.cpp
@@ -56,6 +56,8 @@
     // Some implicit declarations (like intrinsic functions) are not marked.
     // Set the most lenient target on them for maximal flexibility.
     return CFT_HostDevice;
+  } else if (getLangOpts().CUDARelaxedConstexpr && D->isConstexpr()) {
+    return CFT_HostDevice;
   }
 
   return CFT_Host;
Index: test/SemaCUDA/relaxed-constexpr.cu
===================================================================
--- /dev/null
+++ test/SemaCUDA/relaxed-constexpr.cu
@@ -0,0 +1,21 @@
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify -verify-ignore-unexpected=note -DERR %s
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -fcuda-is-device -verify -verify-ignore-unexpected=note -DERR %s
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -fcuda-relaxed-constexpr -verify -verify-ignore-unexpected=note %s
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -fcuda-is-device -fcuda-relaxed-constexpr -verify -verify-ignore-unexpected=note %s
+
+#include "Inputs/cuda.h"
+
+constexpr int foo() { return 0; }
+
+__host__ void host() {
+  constexpr auto x = foo();
+}
+
+__device__ void device() {
+  constexpr auto x = foo();
+#ifdef ERR
+  // expected-error@-2 {{no matching function for call to 'foo'}}
+#else
+  // expected-no-diagnostics
+#endif
+}