Index: include/clang/Sema/Sema.h
===================================================================
--- include/clang/Sema/Sema.h
+++ include/clang/Sema/Sema.h
@@ -3993,11 +3993,15 @@
   // Expression Parsing Callbacks: SemaExpr.cpp.
 
   bool CanUseDecl(NamedDecl *D, bool TreatUnavailableAsInvalid);
+
+  // \param TemplateReceiver is the receiving instantiated template declaration
+  // if it is not a null pointer.
   bool DiagnoseUseOfDecl(NamedDecl *D, ArrayRef<SourceLocation> Locs,
                          const ObjCInterfaceDecl *UnknownObjCClass = nullptr,
                          bool ObjCPropertyAccess = false,
                          bool AvoidPartialAvailabilityChecks = false,
-                         ObjCInterfaceDecl *ClassReciever = nullptr);
+                         ObjCInterfaceDecl *ClassReciever = nullptr,
+                         TemplateDecl *TemplateReceiver = nullptr);
   void NoteDeletedFunction(FunctionDecl *FD);
   void NoteDeletedInheritingConstructor(CXXConstructorDecl *CD);
   std::string getDeletedOrUnavailableSuffix(const FunctionDecl *FD);
@@ -6429,10 +6433,12 @@
 
   bool CheckTemplateArgument(TemplateTypeParmDecl *Param,
                              TypeSourceInfo *Arg);
-  ExprResult CheckTemplateArgument(NonTypeTemplateParmDecl *Param,
-                                   QualType InstantiatedParamType, Expr *Arg,
-                                   TemplateArgument &Converted,
-                               CheckTemplateArgumentKind CTAK = CTAK_Specified);
+  ExprResult
+  CheckTemplateArgument(NonTypeTemplateParmDecl *Param,
+                        QualType InstantiatedParamType, Expr *Arg,
+                        TemplateArgument &Converted,
+                        CheckTemplateArgumentKind CTAK = CTAK_Specified,
+                        TemplateDecl *Template = nullptr);
   bool CheckTemplateTemplateArgument(TemplateParameterList *Params,
                                      TemplateArgumentLoc &Arg);
 
@@ -9975,8 +9981,15 @@
 
 private:
   unsigned ForceCUDAHostDeviceDepth = 0;
+  unsigned IsParsingTemplateArgumentList = 0;
 
 public:
+  void setIsParsingTemplateArgumentList(bool Enable = true) {
+    if (Enable)
+      ++IsParsingTemplateArgumentList;
+    else
+      --IsParsingTemplateArgumentList;
+  }
   /// Increments our count of the number of times we've seen a pragma forcing
   /// functions to be __host__ __device__.  So long as this count is greater
   /// than zero, all functions encountered will be __host__ __device__.
@@ -10196,7 +10209,8 @@
   ///   deferred errors.
   ///
   /// - Otherwise, returns true without emitting any diagnostics.
-  bool CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee);
+  bool CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee,
+                     FunctionDecl *Caller = nullptr);
 
   /// Set __device__ or __host__ __device__ attributes on the given lambda
   /// operator() method.
Index: lib/Parse/ParseTemplate.cpp
===================================================================
--- lib/Parse/ParseTemplate.cpp
+++ lib/Parse/ParseTemplate.cpp
@@ -1296,6 +1296,20 @@
   return Tok.isOneOf(tok::greater, tok::comma);
 }
 
+namespace {
+class ParseTemplateArgumentListRAII {
+  Sema &S;
+
+public:
+  ParseTemplateArgumentListRAII(Sema &_S) : S(_S) {
+    S.setIsParsingTemplateArgumentList();
+  }
+  ~ParseTemplateArgumentListRAII() {
+    S.setIsParsingTemplateArgumentList(false);
+  }
+};
+
+} // namespace
 /// ParseTemplateArgumentList - Parse a C++ template-argument-list
 /// (C++ [temp.names]). Returns true if there was an error.
 ///
@@ -1306,6 +1320,7 @@
 Parser::ParseTemplateArgumentList(TemplateArgList &TemplateArgs) {
 
   ColonProtectionRAIIObject ColonProtection(*this, false);
+  ParseTemplateArgumentListRAII PTAL(Actions);
 
   do {
     ParsedTemplateArgument Arg = ParseTemplateArgument();
Index: lib/Sema/SemaCUDA.cpp
===================================================================
--- lib/Sema/SemaCUDA.cpp
+++ lib/Sema/SemaCUDA.cpp
@@ -833,13 +833,15 @@
   }
 }
 
-bool Sema::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee) {
+bool Sema::CheckCUDACall(SourceLocation Loc, FunctionDecl *Callee,
+                         FunctionDecl *Caller) {
   assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
   assert(Callee && "Callee may not be null.");
   // FIXME: Is bailing out early correct here?  Should we instead assume that
   // the caller is a global initializer?
-  FunctionDecl *Caller = dyn_cast<FunctionDecl>(CurContext);
   if (!Caller)
+    Caller = dyn_cast<FunctionDecl>(CurContext);
+  if (!Caller || IsParsingTemplateArgumentList)
     return true;
 
   // If the caller is known-emitted, mark the callee as known-emitted.
Index: lib/Sema/SemaExpr.cpp
===================================================================
--- lib/Sema/SemaExpr.cpp
+++ lib/Sema/SemaExpr.cpp
@@ -207,7 +207,8 @@
                              const ObjCInterfaceDecl *UnknownObjCClass,
                              bool ObjCPropertyAccess,
                              bool AvoidPartialAvailabilityChecks,
-                             ObjCInterfaceDecl *ClassReceiver) {
+                             ObjCInterfaceDecl *ClassReceiver,
+                             TemplateDecl *Template) {
   SourceLocation Loc = Locs.front();
   if (getLangOpts().CPlusPlus && isa<FunctionDecl>(D)) {
     // If there were any diagnostics suppressed by template argument deduction,
@@ -262,7 +263,11 @@
         DeduceReturnType(FD, Loc))
       return true;
 
-    if (getLangOpts().CUDA && !CheckCUDACall(Loc, FD))
+    if (getLangOpts().CUDA &&
+        !CheckCUDACall(
+            Loc, FD,
+            Template ? dyn_cast<FunctionDecl>(Template->getTemplatedDecl())
+                     : nullptr))
       return true;
   }
 
Index: lib/Sema/SemaTemplate.cpp
===================================================================
--- lib/Sema/SemaTemplate.cpp
+++ lib/Sema/SemaTemplate.cpp
@@ -4753,8 +4753,8 @@
       TemplateArgument Result;
       unsigned CurSFINAEErrors = NumSFINAEErrors;
       ExprResult Res =
-        CheckTemplateArgument(NTTP, NTTPType, Arg.getArgument().getAsExpr(),
-                              Result, CTAK);
+          CheckTemplateArgument(NTTP, NTTPType, Arg.getArgument().getAsExpr(),
+                                Result, CTAK, dyn_cast<TemplateDecl>(Template));
       if (Res.isInvalid())
         return true;
       // If the current template argument causes an error, give up now.
@@ -6123,6 +6123,27 @@
   return true;
 }
 
+namespace {
+bool CheckCUDATemplateArgument(Sema &S, Expr *Arg, TemplateDecl *Template) {
+  if (Template) {
+    Expr *E = Arg;
+    if (UnaryOperator *UO = dyn_cast<UnaryOperator>(E)) {
+      E = UO ? UO->getSubExpr() : nullptr;
+    }
+    if (DeclRefExpr *DRE = dyn_cast_or_null<DeclRefExpr>(E)) {
+      ValueDecl *Entity = DRE ? DRE->getDecl() : nullptr;
+      if (Entity) {
+        if (auto Callee = dyn_cast<FunctionDecl>(Entity))
+          if (auto Caller =
+                  dyn_cast<FunctionDecl>(Template->getTemplatedDecl()))
+            if (!S.CheckCUDACall(Arg->getBeginLoc(), Callee, Caller))
+              return false;
+      }
+    }
+  }
+  return true;
+}
+} // namespace
 /// Check a template argument against its corresponding
 /// non-type template parameter.
 ///
@@ -6133,7 +6154,8 @@
 ExprResult Sema::CheckTemplateArgument(NonTypeTemplateParmDecl *Param,
                                        QualType ParamType, Expr *Arg,
                                        TemplateArgument &Converted,
-                                       CheckTemplateArgumentKind CTAK) {
+                                       CheckTemplateArgumentKind CTAK,
+                                       TemplateDecl *Template) {
   SourceLocation StartLoc = Arg->getBeginLoc();
 
   // If the parameter type somehow involves auto, deduce the type now.
@@ -6530,7 +6552,11 @@
       if (FunctionDecl *Fn = ResolveAddressOfOverloadedFunction(Arg, ParamType,
                                                                 true,
                                                                 FoundResult)) {
-        if (DiagnoseUseOfDecl(Fn, Arg->getBeginLoc()))
+        if (DiagnoseUseOfDecl(Fn, Arg->getBeginLoc(),
+                              /*UnknownObjCClass=*/nullptr,
+                              /*ObjCPropertyAccess=*/false,
+                              /*AvoidPartialAvailabilityChecks=*/false,
+                              /*ClassReciever=*/nullptr, Template))
           return ExprError();
 
         Arg = FixOverloadedFunctionReference(Arg, FoundResult, Fn);
@@ -6539,6 +6565,9 @@
         return ExprError();
     }
 
+    if (!CheckCUDATemplateArgument(*this, Arg, Template))
+      return ExprError();
+
     if (!ParamType->isMemberPointerType()) {
       if (CheckTemplateArgumentAddressOfObjectOrFunction(*this, Param,
                                                          ParamType,
Index: test/SemaCUDA/kernel-template-with-device-func-arg.cu
===================================================================
--- /dev/null
+++ test/SemaCUDA/kernel-template-with-device-func-arg.cu
@@ -0,0 +1,49 @@
+// RUN: %clang_cc1 -fsyntax-only -verify %s
+
+#include "Inputs/cuda.h"
+
+struct C {
+  __device__ void devfun() {}
+  void hostfun() {}
+  template<class T> __device__ void devtempfun() {}
+};
+
+__device__ void devfun() {}
+__host__ void hostfun() {}
+template<class T> __device__ void devtempfun() {}
+
+template <void (*devF)()> __global__ void kernel() { devF();}
+template <typename T, void(T::*devF)()> __global__ void kernel2(T *p) { (p->*devF)(); }
+
+template<> __global__ void kernel<devfun>();
+template<> __global__ void kernel<hostfun>(); // expected-error {{no function template matches function template specialization 'kernel'}}
+                                              // expected-note@-5 {{candidate template ignored: invalid explicitly-specified argument for template parameter 'devF'}}
+template<> __global__ void kernel<devtempfun<int> >();
+
+template<> __global__ void kernel<&devfun>();
+template<> __global__ void kernel<&hostfun>(); // expected-error {{no function template matches function template specialization 'kernel'}}
+                                               // expected-note@-10 {{candidate template ignored: invalid explicitly-specified argument for template parameter 'devF'}}
+template<> __global__ void kernel<&devtempfun<int> >();
+
+template<> __global__ void kernel2<C, &C::devfun>(C *p);
+template<> __global__ void kernel2<C, &C::hostfun>(C *p); // expected-error {{no function template matches function template specialization 'kernel2'}}
+                                                          // expected-note@-14 {{candidate template ignored: invalid explicitly-specified argument for template parameter 'devF'}}
+template<> __global__ void kernel2<C, &C::devtempfun<int> >(C *p);
+
+void fun() {
+  kernel<&devfun><<<1,1>>>();
+  kernel<&hostfun><<<1,1>>>(); // expected-error {{no matching function for call to 'kernel'}}
+                               // expected-note@-21 {{candidate template ignored: invalid explicitly-specified argument for template parameter 'devF'}}
+  kernel<&devtempfun<int> ><<<1,1>>>();
+
+  kernel<devfun><<<1,1>>>();
+  kernel<hostfun><<<1,1>>>(); // expected-error {{no matching function for call to 'kernel'}}
+                              // expected-note@-26 {{candidate template ignored: invalid explicitly-specified argument for template parameter 'devF'}}
+  kernel<devtempfun<int> ><<<1,1>>>();
+
+  C a;
+  kernel2<C, &C::devfun><<<1,1>>>(&a);
+  kernel2<C, &C::hostfun><<<1,1>>>(&a); // expected-error {{no matching function for call to 'kernel2'}}
+                                        // expected-note@-31 {{candidate template ignored: invalid explicitly-specified argument for template parameter 'devF'}}
+  kernel2<C, &C::devtempfun<int> ><<<1,1>>>(&a);
+}