Index: include/clang/Sema/Sema.h
===================================================================
--- include/clang/Sema/Sema.h
+++ include/clang/Sema/Sema.h
@@ -4010,11 +4010,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);
@@ -6453,10 +6457,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);
 
@@ -10224,7 +10230,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/Sema/SemaCUDA.cpp
===================================================================
--- lib/Sema/SemaCUDA.cpp
+++ lib/Sema/SemaCUDA.cpp
@@ -833,12 +833,19 @@
   }
 }
 
-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.");
+
+  if (ExprEvalContexts.back().ExprContext ==
+      ExpressionEvaluationContextRecord::ExpressionKind::EK_TemplateArgument)
+    return true;
+
   // 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)
     return true;
 
Index: lib/Sema/SemaExpr.cpp
===================================================================
--- lib/Sema/SemaExpr.cpp
+++ lib/Sema/SemaExpr.cpp
@@ -213,7 +213,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,
@@ -270,7 +271,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
@@ -4784,8 +4784,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.
@@ -6154,6 +6154,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.
 ///
@@ -6164,7 +6185,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.
@@ -6561,7 +6583,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);
@@ -6570,6 +6596,9 @@
         return ExprError();
     }
 
+    if (!CheckCUDATemplateArgument(*this, Arg, Template))
+      return ExprError();
+
     if (!ParamType->isMemberPointerType()) {
       if (CheckTemplateArgumentAddressOfObjectOrFunction(*this, Param,
                                                          ParamType,