Index: include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- include/clang/Basic/DiagnosticSemaKinds.td
+++ include/clang/Basic/DiagnosticSemaKinds.td
@@ -6067,6 +6067,9 @@
 def err_ref_bad_target : Error<
   "reference to %select{__device__|__global__|__host__|__host__ __device__}0 "
   "function %1 in %select{__device__|__global__|__host__|__host__ __device__}2 function">;
+def warn_host_device_function_calling_host_function : Warning<
+  "calling __host__ function %0 from __host__ __device__ function %1 can lead to runtime errors">,
+  InGroup<CudaCompat>;
 
 def warn_non_pod_vararg_with_format_string : Warning<
   "cannot pass %select{non-POD|non-trivial}0 object of type %1 to variadic "
Index: include/clang/Basic/LangOptions.def
===================================================================
--- include/clang/Basic/LangOptions.def
+++ include/clang/Basic/LangOptions.def
@@ -160,6 +160,7 @@
 LANGOPT(CUDA              , 1, 0, "CUDA")
 LANGOPT(OpenMP            , 1, 0, "OpenMP support")
 LANGOPT(CUDAIsDevice      , 1, 0, "Compiling for CUDA device")
+LANGOPT(CUDAHostDeviceFunctionsCallingHostFunctions, 1, 0, "Allow host device functions to call host functions")
 
 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
@@ -608,6 +608,9 @@
 
 def fcuda_is_device : Flag<["-"], "fcuda-is-device">,
   HelpText<"Generate code for CUDA device">;
+def fcuda_host_device_functions_calling_host_functions : Flag<["-"],
+    "fcuda-host-device-functions-calling-host-functions">,
+  HelpText<"Allow host device functions to call host functions">;
 
 } // let Flags = [CC1Option]
 
Index: lib/Frontend/CompilerInvocation.cpp
===================================================================
--- lib/Frontend/CompilerInvocation.cpp
+++ lib/Frontend/CompilerInvocation.cpp
@@ -253,7 +253,7 @@
     for (unsigned i = 0, e = checkers.size(); i != e; ++i)
       Opts.CheckersControlList.push_back(std::make_pair(checkers[i], enable));
   }
-  
+
   // Go through the analyzer configuration options.
   for (arg_iterator it = Args.filtered_begin(OPT_analyzer_config),
        ie = Args.filtered_end(); it != ie; ++it) {
@@ -1393,6 +1393,9 @@
   if (Args.hasArg(OPT_fcuda_is_device))
     Opts.CUDAIsDevice = 1;
 
+  if (Args.hasArg(OPT_fcuda_host_device_functions_calling_host_functions))
+    Opts.CUDAHostDeviceFunctionsCallingHostFunctions = 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
@@ -92,9 +92,21 @@
     if (Caller->isImplicit()) return false;
 
     bool InDeviceMode = getLangOpts().CUDAIsDevice;
-    if ((InDeviceMode && CalleeTarget != CFT_Device) ||
-        (!InDeviceMode && CalleeTarget != CFT_Host))
+    if (!InDeviceMode && CalleeTarget != CFT_Host)
+        return true;
+    if (InDeviceMode && CalleeTarget != CFT_Device) {
+      // Allow host device functions to call host functions if explicitly
+      // requested.
+      if (CalleeTarget == CFT_Host &&
+          getLangOpts().CUDAHostDeviceFunctionsCallingHostFunctions) {
+        Diag(Caller->getLocation(),
+             diag::warn_host_device_function_calling_host_function)
+            << Callee->getNameAsString() << Caller->getNameAsString();
+        return false;
+      }
+
       return true;
+    }
   }
 
   return false;
Index: test/SemaCUDA/function-target.cu
===================================================================
--- test/SemaCUDA/function-target.cu
+++ test/SemaCUDA/function-target.cu
@@ -1,5 +1,7 @@
-// RUN: %clang_cc1 -fsyntax-only -verify %s
-// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s
+// RUN: %clang_cc1 -fsyntax-only -verify %s -DTEST_HOST
+// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s -DTEST_DEVICE
+// RUN: %clang_cc1 -fsyntax-only -fcuda-host-device-functions-calling-host-functions -verify %s -DTEST_WARN_HD -DTEST_HOST
+// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -fcuda-host-device-functions-calling-host-functions -verify %s -DTEST_WARN_HD -DTEST_DEVICE
 
 #include "Inputs/cuda.h"
 
@@ -32,14 +34,21 @@
   d1g<<<1, 1>>>(); // expected-error {{reference to __global__ function 'd1g' in __device__ function}}
 }
 
-// Expected 0-1 as in one of host/device side compilation it is an error, while
-// not in the other
-__host__ void hd1h(void); // expected-note 0-1 {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
-__device__ void hd1d(void); // expected-note 0-1 {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
+__host__ void hd1h(void);
+#if defined(TEST_DEVICE) && !defined(TEST_WARN_HD)
+// expected-note@-2 {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
+#endif
+__device__ void hd1d(void);
+#ifdef TEST_HOST
+// expected-note@-2 {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
+#endif
 __host__ void hd1hg(void);
 __device__ void hd1dg(void);
 #ifdef __CUDA_ARCH__
-__host__ void hd1hig(void); // expected-note {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
+__host__ void hd1hig(void);
+#if defined(TEST_DEVICE) && !defined(TEST_WARN_HD)
+// expected-note@-2 {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
+#endif
 #else
 __device__ void hd1dig(void); // expected-note {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
 #endif
@@ -47,10 +56,18 @@
 __global__ void hd1g(void); // expected-note {{'hd1g' declared here}}
 
 __host__ __device__ void hd1(void) {
-  // Expected 0-1 as in one of host/device side compilation it is an error,
-  // while not in the other
-  hd1d(); // expected-error 0-1 {{no matching function}}
-  hd1h(); // expected-error 0-1 {{no matching function}}
+#if defined(TEST_WARN_HD) && defined(TEST_DEVICE)
+// expected-warning@-2 {{calling __host__ function hd1h from __host__ __device__ function hd1}}
+// expected-warning@-3 {{calling __host__ function hd1hig from __host__ __device__ function hd1}}
+#endif
+  hd1d();
+#ifdef TEST_HOST
+// expected-error@-2 {{no matching function}}
+#endif
+  hd1h();
+#if defined(TEST_DEVICE) && !defined(TEST_WARN_HD)
+// expected-error@-2 {{no matching function}}
+#endif
 
   // No errors as guarded
 #ifdef __CUDA_ARCH__
@@ -63,8 +80,11 @@
 #ifndef __CUDA_ARCH__
   hd1dig(); // expected-error {{no matching function}}
 #else
-  hd1hig(); // expected-error {{no matching function}}
+  hd1hig();
+#ifndef TEST_WARN_HD
+// expected-error@-2 {{no matching function}}
 #endif
+#endif
 
   hd1hd();
   hd1g<<<1, 1>>>(); // expected-error {{reference to __global__ function 'hd1g' in __host__ __device__ function}}