Index: lib/Sema/SemaCUDA.cpp
===================================================================
--- lib/Sema/SemaCUDA.cpp
+++ lib/Sema/SemaCUDA.cpp
@@ -48,6 +48,12 @@
     if (D->hasAttr<CUDAHostAttr>())
       return CFT_HostDevice;
     return CFT_Device;
+  } else if (D->hasAttr<CUDAHostAttr>()) {
+    return CFT_Host;
+  } else if (D->isImplicit()) {
+    // Some implicit declarations (like intrinsic functions) are not marked.
+    // Set the most lenient target on them for maximal flexibility.
+    return CFT_HostDevice;
   }
 
   return CFT_Host;
Index: test/SemaCUDA/implicit-intrinsic.cu
===================================================================
--- /dev/null
+++ test/SemaCUDA/implicit-intrinsic.cu
@@ -0,0 +1,10 @@
+// RUN: %clang_cc1 -std=gnu++11 -triple nvptx64-unknown-unknown -fsyntax-only -verify %s
+
+#include "Inputs/cuda.h"
+
+// expected-no-diagnostics
+__device__ void __threadfence_system() {
+  // This shouldn't produce an error, since __nvvm_membar_sys is inferred to
+  // be __host__ __device__ and thus callable from device code.
+  __nvvm_membar_sys();
+}