[CUDA] Ignore implicit target attributes during function template instantiation.

[tra]

Summary

Some functions and templates are treated as __host__ __device__ even when they don't have explicitly specified target attributes.
What's worse, this treatment may change depending on command line options (-fno-cuda-host-device-constexpr) or #pragma clang force_cuda_host_device.

Combined with strict checking for matching function target that comes with D25809, it makes it hard to write code which would explicitly instantiate or specialize some functions regardless of pragmas or command line options in effect.

This patch changes the way we match target attributes of base template vs attributes used in explicit instantiation or specialization so that only explicitly specified attributes are considered.

[jlebar]

test/SemaCUDA/function-template-overload.cu

62 ↗

(On Diff #75348)

Oh, I didn't know you could specify a count like this. I have a bunch of tests to fix now!

79 ↗

(On Diff #75348)

D25809's patch description says:

[New behavior]: Require matching target attributes for explicit function template instantiation/specialization and narrow [the] list of candidates to templates with the same target.

Based on this, I do not see why most of the "these should work" candidates work.

Rather than explain here, can you please add a comment to the test?

(At some point someone is probably going to have to write down and make sense of all of the semantics we have invented. It's going to be a pain either way, but at least we can preserve our reasoning in code rather than phab comments.)

[jlebar]

This patch changes the way we match target attributes of base template vs attributes used in explicit instantiation or specialization so that only explicitly specified attributes are considered.

Another question about this: When we have something inside of the force-host-device pragma, the HD attributes it gets are implicit.

This means that the only way to specialize something inside one of these pragmas is to use the pragma around the specialization? That...seems weird?

[tra]

Added a comment explaining expected constexpr function template matching behavior.

[tra]

In D25845#576819, @jlebar wrote:

This patch changes the way we match target attributes of base template vs attributes used in explicit instantiation or specialization so that only explicitly specified attributes are considered.

Another question about this: When we have something inside of the force-host-device pragma, the HD attributes it gets are implicit.

.. and therefore ignored for the purposes of matching specialization to template.

This means that the only way to specialize something inside one of these pragmas is to use the pragma around the specialization? That...seems weird?

The whole purpose of this patch is to *ignore* implicit attributes added by compiler and only consider what user has written in the source code. Specialization will work exactly the same regardless of whether template or its specialization are within the pragma. In all cases we'll only consider attributes explicitly specified by the user. I've used constexpr template to illustrate the case where compiler adds implicit HD attributes to the template, but the logic will apply to the specialization, too -- if attributes are not written, they are not considered when we consider which template to match.

[jlebar]

To close the loop, we talked about this IRL, and I agree this is the most sane option we can come up with. The user-facing principle is that the signatures of function template specializations must match. We consider CUDA attributes to be part of the function's signature. Therefore the target attributes should match.

Now the question is, should the *implicit plus explicit* or *only explicit* target attributes match? To me, implicit+explicit makes more sense on face. But the problem is, we need to allow C++ code like

constexpr template<typename T> void foo(T);
template<> void foo(int);

We implicitly make constexpr functions HD. The template specialization is not constexpr, so is not implicitly HD. But we cannot disallow this code, because it's valid C++. That effectively constrains our choice above to *only explicit* target attributes.

[jlebar]

include/clang/Sema/Sema.h

9396 ↗

(On Diff #75482)

We usually call them "target attributes", not "HD attributes"?

lib/Sema/SemaCUDA.cpp

97 ↗

(On Diff #75482)

hasAttr?

900 ↗

(On Diff #75482)

I don't see where this function is declared?

lib/Sema/SemaDeclAttr.cpp

5628 ↗

(On Diff #75482)

Is this a functional change in this patch? We don't check for duplicates of most other attributes, so I'm not sure why it should matter with these. If it does matter, we should have comments...

lib/Sema/SemaTemplate.cpp

7046 ↗

(On Diff #75482)

http://jlebar.com/2011/12/16/Boolean_parameters_to_API_functions_considered_harmful..html :)

At least please do

IdentifyCUDATarget(Specialization, /* IgnoreImplicitTargetAttrs = */ true);

7168 ↗

(On Diff #75482)

I am unsure of whether or why we need this anymore, since I don't see this declared in any header, and also since I thought we weren't merging attrs anymore?

8119 ↗

(On Diff #75482)

Same here wrt bool param.

[tra]

addressed jlebar's comments.

[tra]

include/clang/Sema/Sema.h

9396 ↗

(On Diff #75482)

Implementation only ignores H and D attributes. Global is never explicit (AFAICT) and we never ignore InvalidTarget even if it's implicit.

lib/Sema/SemaDeclAttr.cpp

5628 ↗

(On Diff #75482)

This function copies over attributes from template to instantiation, so when we already had explicit attributes we would end up with another one which has potential to mess with our attempt to ignore implicit attributes. getAttr() would return the first attribute it finds and if it happens to be implicit one, explicit one would be ignored.

Plus, it was rather weird to see duplicate attributes hanging off FunctionDecls in AST.

lib/Sema/SemaTemplate.cpp

7168 ↗

(On Diff #75482)

It's gone as it's not needed any more. It used to be needed when I disabled copying of target attributes from the base template.

[jlebar]

lib/Sema/SemaDeclAttr.cpp

5628 ↗

(On Diff #75482)

This function copies over attributes from template to instantiation

It does? I thought this just handled seeing an attribute in the source?

I am all for having a pretty AST, but

• if the user explicitly puts two __host__ attributes on the function, clearly we don't care about having the attribute twice on the function, and alternatively
• if the user explicitly puts a __host__ attribute on the function and we also somehow got an implicit attr there, this check seems wrong, since it will let the implicit one win over the explicit one?

lib/Sema/SemaTemplate.cpp

7046 ↗

(On Diff #75482)

Please put the equals sign. Without it, who knows whether the following call is sync or async:

make_call(/* async */ false);

Also using the equals sign lets a clang-tidy check ensure that we're using the right name.

8119 ↗

(On Diff #75482)

Same here.

[tra]

• Instead of relying on the first attribute we find, check all matching ones.
• Specializations inherit their target attributes from their base template only. Their effective target always matches that of the template and is no longer affected by whether specialization differens from template in its constexpr-ness.

[jlebar]

Doesn't look like we changed any testcases when we changed this behavior? The change in behavior may be unobservable, but even still it seems worthwhile to have tests that check that we're not doing any of the alternatives we considered.

lib/Sema/SemaCUDA.cpp

99 ↗

(On Diff #75812)

Is this early return necessary?

107 ↗

(On Diff #75812)

Could we write this function as

return llvm::any_of(D->getAttrs(), [&](Attr *Attribute) {
  return isa<A>(Attribute) && (!IgnoreImplicitAttr || !Attribute->isImplicit()) });

893 ↗

(On Diff #75812)

Perhaps ToFD, FromFD would be better names than "new" and "old".

896 ↗

(On Diff #75812)

This comment is confusing because this function's implementation doesn't have anything to do with templates, but the first sentence says we're copying from a template to "FD".

In addition to cleaning this up, maybe we should move the comment into Sema.h?

912 ↗

(On Diff #75812)

Could this be written as

template<typename AttrTy>
static void copyAttrIfPresent(FunctionDecl *NewFD, FunctionDecl *OldFD) {
  if (A *Attribute = OldFD->getAttr<A>()) {
    ...
  }
}

void Sema::inheritCUDATargetAttrs(...) {
  copyAttrIfPresent<CUDAGlobalAttr>(NewFD, OldFD);
  ...
}

? In particular I am not sure why we need this particular control flow in this function.

lib/Sema/SemaDecl.cpp

8372 ↗

(On Diff #75812)

"from their templates" "in *the* foo() call below".

Although because "from their templates" is kind of confusing, I might rewrite to be singular:

Instead, a specialization inherits its target attributes from its template ...

lib/Sema/SemaTemplate.cpp

7048 ↗

(On Diff #75812)

Missing some articles:

Target attributes are part of the cuda function signature, so the deduced template's cuda target must match XXX [1].  Given that regular template deduction [2] does not take target attributes into account, we reject candidates here that have a different target.

[1] I am not sure what XXX should be. The deduced template's cuda target must match what, exactly?

[2] What is "regular template deduction"?

7171 ↗

(On Diff #75812)

Suggest

A function template specialization inherits the target attributes of its template.  (We require the attributes explicitly in the code to match, but a template may have implicit attributes by virtue e.g. of being constexpr, and it passes these implicit attributes on to its specializations.)

7173 ↗

(On Diff #75812)

Isn't this backwards? The function is to, from? If this is a bug, can we add a test to catch this?

[tra]

Addressed Justin's comments.

lib/Sema/SemaCUDA.cpp

99 ↗

(On Diff #75812)

Yes. Otherwise D->getAttrs() will trigger assert(hasAttrs) if we don't have any attributes.

107 ↗

(On Diff #75812)

Done.

896 ↗

(On Diff #75812)

I've changed second argument type to const FunctionTemplateDecl* to reflect my intent a bit better.
Moved the comment to Sema.h.

lib/Sema/SemaTemplate.cpp

7048 ↗

(On Diff #75812)

[1] it must match the target of its template.
[2] "C++ template deduction"
Rephrased the comment based on IRL conversation w/ jlebar@.

7173 ↗

(On Diff #75812)

The arguments have different types now, one of them const&, so it should be unambiguous what's input and what's output.
As for order of inputs/outputs, my mental model is result=input as in memcpy(dest, src). Style guide does not seem to say anything on order of input/output arguments.

[jlebar]

lib/Sema/SemaDecl.cpp

8415 ↗

(On Diff #80677)

"the Foo() call below", or "in Foo() below". (I guess "Foo()" is like a proper name. :)

lib/Sema/SemaTemplate.cpp

7048 ↗

(On Diff #75812)

It's not showing a rephrased comment above; maybe you forgot to commit before arc diff?

test/SemaCUDA/function-template-overload.cu

73 ↗

(On Diff #80677)

"the implicit HD attributes the compiler gives" :)

74 ↗

(On Diff #80677)

comma before "so" (conjunction joining independent clauses)

[tra]

Fixed comments.

[tra]

lib/Sema/SemaTemplate.cpp

7048 ↗

(On Diff #75812)

Argh. I've apparently deleted corrected version after I've showed it to you. :-/ Fixed now.

[jlebar]

lgtm!

[RKSimon]

cfe/trunk/lib/Sema/SemaDecl.cpp

8416

@tra Sorry for touching such an old ticket - but cppcheck is warning that we are using a boolean result in a bitwise expression - I'm assuming this should be:

if (getLangOpts().CUDA && !isFunctionTemplateSpecialization)

[tra]

cfe/trunk/lib/Sema/SemaDecl.cpp

8416

Good catch. I'll fix it.

[tra]

cfe/trunk/lib/Sema/SemaDecl.cpp

8416

Fixed in rL369777.

cfe/trunk/include/clang/Sema/Sema.h

enum CUDAFunctionTarget {
     CFT_InvalidTarget
   };
 
   /// Determines whether the given function is a CUDA device/host/kernel/etc.
   /// function.
   ///
   /// Use this rather than examining the function's attributes yourself -- you
   /// will get it wrong.  Returns CFT_Host if D is null.
-  CUDAFunctionTarget IdentifyCUDATarget(const FunctionDecl *D);
+  CUDAFunctionTarget IdentifyCUDATarget(const FunctionDecl *D,
+                                        bool IgnoreImplicitHDAttr = false);
   CUDAFunctionTarget IdentifyCUDATarget(const AttributeList *Attr);
 
   /// Gets the CUDA target for the current context.
   CUDAFunctionTarget CurrentCUDATarget() {
     return IdentifyCUDATarget(dyn_cast<FunctionDecl>(CurContext));
   }
 
   // CUDA function call preference. Must be ordered numerically from
public:
 
   /// \return true if \p CD can be considered empty according to CUDA
   /// (E.2.3.1 in CUDA 7.5 Programming guide).
   bool isEmptyCudaConstructor(SourceLocation Loc, CXXConstructorDecl *CD);
   bool isEmptyCudaDestructor(SourceLocation Loc, CXXDestructorDecl *CD);
 
   /// Check whether NewFD is a valid overload for CUDA. Emits
   /// diagnostics and invalidates NewFD if not.
-  void checkCUDATargetOverload(FunctionDecl *NewFD, LookupResult &Previous);
+  void checkCUDATargetOverload(FunctionDecl *NewFD,
+                               const LookupResult &Previous);
+  /// Copies target attributes from the template TD to the function FD.
+  void inheritCUDATargetAttrs(FunctionDecl *FD, const FunctionTemplateDecl &TD);
 
   /// \name Code completion
   //@{
   /// \brief Describes the context in which code completion occurs.
   enum ParserCompletionContext {
     /// \brief Code completion occurs at top-level or namespace context.
     PCC_Namespace,
     /// \brief Code completion occurs within a class, struct, or union.

cfe/trunk/lib/Sema/SemaCUDA.cpp

if (HasHostAttr && HasDeviceAttr)
     return CFT_HostDevice;
 
   if (HasDeviceAttr)
     return CFT_Device;
 
   return CFT_Host;
 }
 
+template <typename A>
+static bool hasAttr(const FunctionDecl *D, bool IgnoreImplicitAttr) {
+  return D->hasAttrs() && llvm::any_of(D->getAttrs(), [&](Attr *Attribute) {
+           return isa<A>(Attribute) &&
+                  !(IgnoreImplicitAttr && Attribute->isImplicit());
+         });
+}
+
 /// IdentifyCUDATarget - Determine the CUDA compilation target for this function
-Sema::CUDAFunctionTarget Sema::IdentifyCUDATarget(const FunctionDecl *D) {
+Sema::CUDAFunctionTarget Sema::IdentifyCUDATarget(const FunctionDecl *D,
+                                                  bool IgnoreImplicitHDAttr) {
   // Code that lives outside a function is run on the host.
   if (D == nullptr)
     return CFT_Host;
 
   if (D->hasAttr<CUDAInvalidTargetAttr>())
     return CFT_InvalidTarget;
 
   if (D->hasAttr<CUDAGlobalAttr>())
     return CFT_Global;
 
-  if (D->hasAttr<CUDADeviceAttr>()) {
-    if (D->hasAttr<CUDAHostAttr>())
+  if (hasAttr<CUDADeviceAttr>(D, IgnoreImplicitHDAttr)) {
+    if (hasAttr<CUDAHostAttr>(D, IgnoreImplicitHDAttr))
       return CFT_HostDevice;
     return CFT_Device;
-  } else if (D->hasAttr<CUDAHostAttr>()) {
+  } else if (hasAttr<CUDAHostAttr>(D, IgnoreImplicitHDAttr)) {
     return CFT_Host;
-  } else if (D->isImplicit()) {
+  } else if (D->isImplicit() && !IgnoreImplicitHDAttr) {
     // 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;
 }
 
if (Target == CFT_Global || Target == CFT_Device) {
     Method->addAttr(CUDADeviceAttr::CreateImplicit(Context));
   } else if (Target == CFT_HostDevice) {
     Method->addAttr(CUDADeviceAttr::CreateImplicit(Context));
     Method->addAttr(CUDAHostAttr::CreateImplicit(Context));
   }
 }
 
 void Sema::checkCUDATargetOverload(FunctionDecl *NewFD,
-                                   LookupResult &Previous) {
+                                   const LookupResult &Previous) {
   assert(getLangOpts().CUDA && "Should only be called during CUDA compilation");
   CUDAFunctionTarget NewTarget = IdentifyCUDATarget(NewFD);
   for (NamedDecl *OldND : Previous) {
     FunctionDecl *OldFD = OldND->getAsFunction();
     if (!OldFD)
       continue;
 
     CUDAFunctionTarget OldTarget = IdentifyCUDATarget(OldFD);
if (NewTarget != OldTarget &&
       Diag(NewFD->getLocation(), diag::err_cuda_ovl_target)
           << NewTarget << NewFD->getDeclName() << OldTarget << OldFD;
       Diag(OldFD->getLocation(), diag::note_previous_declaration);
       NewFD->setInvalidDecl();
       break;
     }
   }
 }
+
+template <typename AttrTy>
+static void copyAttrIfPresent(Sema &S, FunctionDecl *FD,
+                              const FunctionDecl &TemplateFD) {
+  if (AttrTy *Attribute = TemplateFD.getAttr<AttrTy>()) {
+    AttrTy *Clone = Attribute->clone(S.Context);
+    Clone->setInherited(true);
+    FD->addAttr(Clone);
+  }
+}
+
+void Sema::inheritCUDATargetAttrs(FunctionDecl *FD,
+                                  const FunctionTemplateDecl &TD) {
+  const FunctionDecl &TemplateFD = *TD.getTemplatedDecl();
+  copyAttrIfPresent<CUDAGlobalAttr>(*this, FD, TemplateFD);
+  copyAttrIfPresent<CUDAHostAttr>(*this, FD, TemplateFD);
+  copyAttrIfPresent<CUDADeviceAttr>(*this, FD, TemplateFD);
+}

cfe/trunk/lib/Sema/SemaDecl.cpp

if (UnifySection(CodeSegStack.CurrentValue->getString(),
                          ASTContext::PSF_Read,
                      NewFD))
       NewFD->dropAttr<SectionAttr>();
   }
 
   // Handle attributes.
   ProcessDeclAttributes(S, NewFD, D);
 
-  if (getLangOpts().CUDA)
-    maybeAddCUDAHostDeviceAttrs(NewFD, Previous);
-
   if (getLangOpts().OpenCL) {
     // OpenCL v1.1 s6.5: Using an address space qualifier in a function return
     // type declaration will generate a compilation error.
     unsigned AddressSpace = NewFD->getReturnType().getAddressSpace();
     if (AddressSpace == LangAS::opencl_local ||
         AddressSpace == LangAS::opencl_global ||
         AddressSpace == LangAS::opencl_constant) {
       Diag(NewFD->getLocation(),
if (D.getName().getKind() == UnqualifiedId::IK_TemplateId) {
       // which we're recovering from as if the user had written:
       //   friend void foo<>(int);
       // Go ahead and fake up a template id.
       HasExplicitTemplateArgs = true;
       TemplateArgs.setLAngleLoc(D.getIdentifierLoc());
       TemplateArgs.setRAngleLoc(D.getIdentifierLoc());
     }
 
+    // We do not add HD attributes to specializations here because
+    // they may have different constexpr-ness compared to their
+    // templates and, after maybeAddCUDAHostDeviceAttrs() is applied,
+    // may end up with different effective targets. Instead, a
+    // specialization inherits its target attributes from its template
+    // in the CheckFunctionTemplateSpecialization() call below.
+    if (getLangOpts().CUDA & !isFunctionTemplateSpecialization)
+      maybeAddCUDAHostDeviceAttrs(NewFD, Previous);
+
     // If it's a friend (and only if it's a friend), it's possible
     // that either the specialized function type or the specialized
     // template is dependent, and therefore matching will fail.  In
     // this case, don't check the specialization yet.
     bool InstantiationDependent = false;
     if (isFunctionTemplateSpecialization && isFriend &&
         (NewFD->getType()->isDependentType() || DC->isDependentContext() ||
          TemplateSpecializationType::anyDependentTemplateArguments(

cfe/trunk/lib/Sema/SemaTemplate.cpp

if (FunctionTemplateDecl *FunTmpl = dyn_cast<FunctionTemplateDecl>(Ovl)) {
         // that we can provide nifty diagnostics.
         FailedCandidates.addCandidate().set(
             I.getPair(), FunTmpl->getTemplatedDecl(),
             MakeDeductionFailureInfo(Context, TDK, Info));
         (void)TDK;
         continue;
       }
 
-      // Target attributes are part of function signature during cuda
-      // compilation, so deduced template must also have matching CUDA
-      // target. Given that regular template deduction does not take
-      // target attributes into account, we perform target match check
-      // here and reject candidates that have different target.
+      // Target attributes are part of the cuda function signature, so
+      // the deduced template's cuda target must match that of the
+      // specialization.  Given that C++ template deduction does not
+      // take target attributes into account, we reject candidates
+      // here that have a different target.
       if (LangOpts.CUDA &&
-          IdentifyCUDATarget(Specialization) != IdentifyCUDATarget(FD)) {
+          IdentifyCUDATarget(Specialization,
+                             /* IgnoreImplicitHDAttributes = */ true) !=
+              IdentifyCUDATarget(FD, /* IgnoreImplicitHDAttributes = */ true)) {
         FailedCandidates.addCandidate().set(
             I.getPair(), FunTmpl->getTemplatedDecl(),
             MakeDeductionFailureInfo(Context, TDK_CUDATargetMismatch, Info));
         continue;
       }
 
       // Record this candidate.
       if (ExplicitTemplateArgs)
bool Sema::CheckFunctionTemplateSpecialization(
   // Take copies of (semantic and syntactic) template argument lists.
   const TemplateArgumentList* TemplArgs = new (Context)
     TemplateArgumentList(Specialization->getTemplateSpecializationArgs());
   FD->setFunctionTemplateSpecialization(
       Specialization->getPrimaryTemplate(), TemplArgs, /*InsertPos=*/nullptr,
       SpecInfo->getTemplateSpecializationKind(),
       ExplicitTemplateArgs ? &ConvertedTemplateArgs[Specialization] : nullptr);
 
+  // A function template specialization inherits the target attributes
+  // of its template.  (We require the attributes explicitly in the
+  // code to match, but a template may have implicit attributes by
+  // virtue e.g. of being constexpr, and it passes these implicit
+  // attributes on to its specializations.)
+  if (LangOpts.CUDA)
+    inheritCUDATargetAttrs(FD, *Specialization->getPrimaryTemplate());
+
   // The "previous declaration" for this function template specialization is
   // the prior function template specialization.
   Previous.clear();
   Previous.addDecl(Specialization);
   return false;
 }
 
 /// \brief Perform semantic analysis for the given non-template member
if (TemplateDeductionResult TDK
       // Keep track of almost-matches.
       FailedCandidates.addCandidate()
           .set(P.getPair(), FunTmpl->getTemplatedDecl(),
                MakeDeductionFailureInfo(Context, TDK, Info));
       (void)TDK;
       continue;
     }
 
-    // Target attributes are part of function signature during cuda
-    // compilation, so deduced template must also have matching CUDA
-    // target. Given that regular template deduction does not take it
-    // into account, we perform target match check here and reject
-    // candidates that have different target.
-    if (LangOpts.CUDA) {
-      CUDAFunctionTarget DeclaratorTarget = IdentifyCUDATarget(Attr);
-      // We need to adjust target when HD is forced by
-      // #pragma clang force_cuda_host_device
-      if (ForceCUDAHostDeviceDepth > 0 &&
-          (DeclaratorTarget == CFT_Device || DeclaratorTarget == CFT_Host))
-        DeclaratorTarget = CFT_HostDevice;
-      if (IdentifyCUDATarget(Specialization) != DeclaratorTarget) {
-        FailedCandidates.addCandidate().set(
-            P.getPair(), FunTmpl->getTemplatedDecl(),
-            MakeDeductionFailureInfo(Context, TDK_CUDATargetMismatch, Info));
-        continue;
-      }
-    }
+    // Target attributes are part of the cuda function signature, so
+    // the cuda target of the instantiated function must match that of its
+    // template.  Given that C++ template deduction does not take
+    // target attributes into account, we reject candidates here that
+    // have a different target.
+    if (LangOpts.CUDA &&
+        IdentifyCUDATarget(Specialization,
+                           /* IgnoreImplicitHDAttributes = */ true) !=
+            IdentifyCUDATarget(Attr)) {
 
     Matches.addDecl(Specialization, P.getAccess());
   }
 
   // Find the most specialized function template specialization.
   UnresolvedSetIterator Result = getMostSpecialized(
       Matches.begin(), Matches.end(), FailedCandidates,
       D.getIdentifierLoc(),

cfe/trunk/test/SemaCUDA/function-template-overload.cu

 template  <> __device__ __host__ DType overload_h_d(long a); // There's no HD template...
 // expected-error@-1 {{no function template matches function template specialization 'overload_h_d'}}
 template  <> __device__ __host__ HType overload_h_d(long a); // There's no HD template...
 // expected-error@-1 {{no function template matches function template specialization 'overload_h_d'}}
 template  <> __device__ DType overload_h_d(long a); // OK. instantiates D
 template  <> __host__ HType overload_h_d(long a); // OK. instantiates H
 
 
-// Can't overload HD template with H or D template, though functions are OK.
+// Can't overload HD template with H or D template, though
+// non-template functions are OK.
 template <typename T> __host__ __device__ HDType overload_hd(T a) { return HDType(); }
 // expected-note@-1 {{previous declaration is here}}
 // expected-note@-2 2 {{candidate template ignored: could not match 'HDType' against 'HType'}}
 template <typename T> __device__ HDType overload_hd(T a);
 // expected-error@-1 {{__device__ function 'overload_hd' cannot overload __host__ __device__ function 'overload_hd'}}
 __device__ HDType overload_hd(int a); // OK.
 
 // Verify that target attributes are taken into account when we
 // explicitly specialize or instantiate function tempaltes.
 template <> __host__ HType overload_hd(int a);
 // expected-error@-1 {{no function template matches function template specialization 'overload_hd'}}
 template __host__ HType overload_hd(long a);
 // expected-error@-1 {{explicit instantiation of 'overload_hd' does not refer to a function template, variable template, member function, member class, or static data member}}
 __host__ HType overload_hd(int a); // OK
 
 template <typename T> __host__ T overload_h(T a); // expected-note {{previous declaration is here}}
 template <typename T> __host__ __device__ T overload_h(T a);
 // expected-error@-1 {{__host__ __device__ function 'overload_h' cannot overload __host__ function 'overload_h'}}
 template <typename T> __device__ T overload_h(T a); // OK. D can overload H.
 
 template <typename T> __host__ HType overload_h_d2(T a) { return HType(); }
 template <typename T> __host__ __device__ HDType overload_h_d2(T a) { return HDType(); }
 template <typename T1, typename T2 = int> __device__ DType overload_h_d2(T1 a) { T1 x; T2 y; return DType(); }
 
+// constexpr functions are implicitly HD, but explicit
+// instantiation/specialization must use target attributes as written.
+template <typename T> constexpr T overload_ce_implicit_hd(T a) { return a+1; }
+// expected-note@-1 3 {{candidate template ignored: target attributes do not match}}
+
+// These will not match the template.
+template __host__ __device__ int overload_ce_implicit_hd(int a);
+// expected-error@-1 {{explicit instantiation of 'overload_ce_implicit_hd' does not refer to a function template, variable template, member function, member class, or static data member}}
+template <> __host__ __device__ long overload_ce_implicit_hd(long a);
+// expected-error@-1 {{no function template matches function template specialization 'overload_ce_implicit_hd'}}
+template <> __host__ __device__ constexpr long overload_ce_implicit_hd(long a);
+// expected-error@-1 {{no function template matches function template specialization 'overload_ce_implicit_hd'}}
+
+// These should work, because template matching ignores the implicit
+// HD attributes the compiler gives to constexpr functions/templates,
+// so 'overload_ce_implicit_hd' template will match __host__ functions
+// only.
+template __host__ int overload_ce_implicit_hd(int a);
+template <> __host__ long overload_ce_implicit_hd(long a);
+
+template float overload_ce_implicit_hd(float a);
+template <> float* overload_ce_implicit_hd(float *a);
+template <> constexpr double overload_ce_implicit_hd(double a) { return a + 3.0; };
+
 __host__ void hf() {
   overload_hd(13);
+  overload_ce_implicit_hd('h');        // Implicitly instantiated
+  overload_ce_implicit_hd(1.0f);       // Explicitly instantiated
+  overload_ce_implicit_hd(2.0);        // Explicitly specialized
 
   HType h = overload_h_d(10);
   HType h2i = overload_h_d2<int>(11);
   HType h2ii = overload_h_d2<int>(12);
 
   // These should be implicitly instantiated from __host__ template returning HType.
-  DType d = overload_h_d(20); // expected-error {{no viable conversion from 'HType' to 'DType'}}
-  DType d2i = overload_h_d2<int>(21); // expected-error {{no viable conversion from 'HType' to 'DType'}}
+  DType d = overload_h_d(20);          // expected-error {{no viable conversion from 'HType' to 'DType'}}
+  DType d2i = overload_h_d2<int>(21);  // expected-error {{no viable conversion from 'HType' to 'DType'}}
   DType d2ii = overload_h_d2<int>(22); // expected-error {{no viable conversion from 'HType' to 'DType'}}
 }
 __device__ void df() {
   overload_hd(23);
+  overload_ce_implicit_hd('d');        // Implicitly instantiated
+  overload_ce_implicit_hd(1.0f);       // Explicitly instantiated
+  overload_ce_implicit_hd(2.0);        // Explicitly specialized
 
   // These should be implicitly instantiated from __device__ template returning DType.
-  HType h = overload_h_d(10); // expected-error {{no viable conversion from 'DType' to 'HType'}}
-  HType h2i = overload_h_d2<int>(11); // expected-error {{no viable conversion from 'DType' to 'HType'}}
+  HType h = overload_h_d(10);          // expected-error {{no viable conversion from 'DType' to 'HType'}}
+  HType h2i = overload_h_d2<int>(11);  // expected-error {{no viable conversion from 'DType' to 'HType'}}
   HType h2ii = overload_h_d2<int>(12); // expected-error {{no viable conversion from 'DType' to 'HType'}}
 
   DType d = overload_h_d(20);
   DType d2i = overload_h_d2<int>(21);
   DType d2ii = overload_h_d2<int>(22);
 }