[Clang][HIP] Target-dependent overload resolution in declarators and specifiers

So far, the resolution of host/device overloads for functions in HIP/CUDA
operates as if in a host-device context for code outside of function bodies,
e.g., in expressions that are part of template arguments in top-level
declarations. This means that, if separate host and device overloads are
declared, the device overload is used in the device compilation phase and the
host overload is used in the host compilation phase.

This patch changes overload resolution in such cases to prefer overloads that
match the target of the declaration in which they occur. For example:

__device__ constexpr int get_n() { return 64; }
__host__ constexpr int get_n() { return -1; }
__device__ std::enable_if<(get_n() > 32)>::type foo() { }

Before, this code would not compile, because get_n resolved to the host
overload during host compilation, causing an error. With this patch, the call
to get_n in the declaration of the device function foo resolves to the device
overload in host and device compilation.

If attributes that affect the declaration's target occur after a call with
target-dependent overload resolution, a warning is issued. This is realized by
registering the Kinds of relevant attributes in the CUDATargetContext when they
are parsed.

This is an alternative to PR llvm#93546, which is required for PR llvm#91478.

Author: ritter-x2a

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -9017,6 +9017,10 @@ def err_global_call_not_config : Error<
 def err_ref_bad_target : Error<
   "reference to %select{__device__|__global__|__host__|__host__ __device__}0 "
   "%select{function|variable}1 %2 in %select{__device__|__global__|__host__|__host__ __device__}3 function">;
+def warn_target_specfier_ignored : Warning<
+  "target specifier has been ignored for overload resolution; "
+  "move the target specifier to the beginning of the declaration to use it for overload resolution">,
+  InGroup<IgnoredAttributes>;
 def note_cuda_const_var_unpromoted : Note<
   "const variable cannot be emitted on device side due to dynamic initialization">;
 def note_cuda_host_var : Note<

clang/include/clang/Sema/SemaCUDA.h

@@ -104,6 +104,8 @@ class SemaCUDA : public SemaBase {
   CUDAFunctionTarget IdentifyTarget(const FunctionDecl *D,
                                     bool IgnoreImplicitHDAttr = false);
   CUDAFunctionTarget IdentifyTarget(const ParsedAttributesView &Attrs);
+  CUDAFunctionTarget IdentifyTarget(
+      const SmallVectorImpl<clang::AttributeCommonInfo::Kind> &AttrKinds);
 
   enum CUDAVariableTarget {
     CVT_Device,  /// Emitted on device side with a shadow variable on host side
@@ -120,21 +122,43 @@ class SemaCUDA : public SemaBase {
     CTCK_Unknown,       /// Unknown context
     CTCK_InitGlobalVar, /// Function called during global variable
                         /// initialization
+    CTCK_Declaration,   /// Function called in a declaration specifier or
+                        /// declarator outside of other contexts, usually in
+                        /// template arguments.
   };
 
   /// Define the current global CUDA host/device context where a function may be
   /// called. Only used when a function is called outside of any functions.
-  struct CUDATargetContext {
-    CUDAFunctionTarget Target = CUDAFunctionTarget::HostDevice;
+  class CUDATargetContext {
+  public:
     CUDATargetContextKind Kind = CTCK_Unknown;
-    Decl *D = nullptr;
+
+    CUDATargetContext() = default;
+
+    CUDATargetContext(SemaCUDA *S, CUDATargetContextKind Kind,
+                      CUDAFunctionTarget Target);
+
+    CUDAFunctionTarget getTarget();
+
+    /// If this is a CTCK_Declaration context, update the Target based on Attrs.
+    /// No-op otherwise.
+    /// Issues a diagnostic if the target changes after it has been queried
+    /// before.
+    void tryRegisterTargetAttrs(const ParsedAttributesView &Attrs);
+
+  private:
+    SemaCUDA *S = nullptr;
+    CUDAFunctionTarget Target = CUDAFunctionTarget::HostDevice;
+    SmallVector<clang::AttributeCommonInfo::Kind, 0> AttrKinds;
+    bool TargetQueried = false;
+
   } CurCUDATargetCtx;
 
   struct CUDATargetContextRAII {
     SemaCUDA &S;
     SemaCUDA::CUDATargetContext SavedCtx;
     CUDATargetContextRAII(SemaCUDA &S_, SemaCUDA::CUDATargetContextKind K,
-                          Decl *D);
+                          Decl *D = nullptr);
     ~CUDATargetContextRAII() { S.CurCUDATargetCtx = SavedCtx; }
   };
 

clang/lib/Parse/ParseDecl.cpp

@@ -311,6 +311,9 @@ void Parser::ParseGNUAttributes(ParsedAttributes &Attrs,
   }
 
   Attrs.Range = SourceRange(StartLoc, EndLoc);
+
+  if (Actions.getLangOpts().CUDA)
+    Actions.CUDA().CurCUDATargetCtx.tryRegisterTargetAttrs(Attrs);
 }
 
 /// Determine whether the given attribute has an identifier argument.
@@ -1003,6 +1006,9 @@ void Parser::ParseMicrosoftDeclSpecs(ParsedAttributes &Attrs) {
   }
 
   Attrs.Range = SourceRange(StartLoc, EndLoc);
+
+  if (Actions.getLangOpts().CUDA)
+    Actions.CUDA().CurCUDATargetCtx.tryRegisterTargetAttrs(Attrs);
 }
 
 void Parser::ParseMicrosoftTypeAttributes(ParsedAttributes &attrs) {

clang/lib/Parse/ParseDeclCXX.cpp

@@ -27,6 +27,7 @@
 #include "clang/Sema/EnterExpressionEvaluationContext.h"
 #include "clang/Sema/ParsedTemplate.h"
 #include "clang/Sema/Scope.h"
+#include "clang/Sema/SemaCUDA.h"
 #include "clang/Sema/SemaCodeCompletion.h"
 #include "llvm/ADT/SmallString.h"
 #include "llvm/Support/TimeProfiler.h"
@@ -2852,6 +2853,11 @@ Parser::DeclGroupPtrTy Parser::ParseCXXClassMemberDeclaration(
     ParsedTemplateInfo &TemplateInfo, ParsingDeclRAIIObject *TemplateDiags) {
   assert(getLangOpts().CPlusPlus &&
          "ParseCXXClassMemberDeclaration should only be called in C++ mode");
+  SemaCUDA::CUDATargetContextRAII CTCRAII(Actions.CUDA(),
+                                          SemaCUDA::CTCK_Declaration);
+  if (Actions.getLangOpts().CUDA)
+    Actions.CUDA().CurCUDATargetCtx.tryRegisterTargetAttrs(AccessAttrs);
+
   if (Tok.is(tok::at)) {
     if (getLangOpts().ObjC && NextToken().isObjCAtKeyword(tok::objc_defs))
       Diag(Tok, diag::err_at_defs_cxx);

clang/lib/Parse/Parser.cpp

@@ -21,6 +21,7 @@
 #include "clang/Sema/DeclSpec.h"
 #include "clang/Sema/ParsedTemplate.h"
 #include "clang/Sema/Scope.h"
+#include "clang/Sema/SemaCUDA.h"
 #include "clang/Sema/SemaCodeCompletion.h"
 #include "llvm/Support/Path.h"
 #include "llvm/Support/TimeProfiler.h"
@@ -1133,6 +1134,13 @@ bool Parser::isStartOfFunctionDefinition(const ParsingDeclarator &Declarator) {
 Parser::DeclGroupPtrTy Parser::ParseDeclOrFunctionDefInternal(
     ParsedAttributes &Attrs, ParsedAttributes &DeclSpecAttrs,
     ParsingDeclSpec &DS, AccessSpecifier AS) {
+  SemaCUDA::CUDATargetContextRAII CTCRAII(Actions.CUDA(),
+                                          SemaCUDA::CTCK_Declaration);
+  if (Actions.getLangOpts().CUDA) {
+    Actions.CUDA().CurCUDATargetCtx.tryRegisterTargetAttrs(Attrs);
+    Actions.CUDA().CurCUDATargetCtx.tryRegisterTargetAttrs(DeclSpecAttrs);
+  }
+
   // Because we assume that the DeclSpec has not yet been initialised, we simply
   // overwrite the source range and attribute the provided leading declspec
   // attributes.

clang/lib/Sema/SemaCUDA.cpp

@@ -18,6 +18,7 @@
 #include "clang/Basic/TargetInfo.h"
 #include "clang/Lex/Preprocessor.h"
 #include "clang/Sema/Lookup.h"
+#include "clang/Sema/ParsedAttr.h"
 #include "clang/Sema/ScopeInfo.h"
 #include "clang/Sema/Sema.h"
 #include "clang/Sema/SemaDiagnostic.h"
@@ -68,13 +69,28 @@ ExprResult SemaCUDA::ActOnExecConfigExpr(Scope *S, SourceLocation LLLLoc,
                                /*IsExecConfig=*/true);
 }
 
-CUDAFunctionTarget SemaCUDA::IdentifyTarget(const ParsedAttributesView &Attrs) {
+namespace {
+
+// This iterator adaptor enables sharing a IdentifyTarget implementation for
+// ParsedAttributesView and for vectors of AttributeCommonInfo::Kind.
+struct AttrKindIterator
+    : llvm::iterator_adaptor_base<
+          AttrKindIterator, ParsedAttributesView::const_iterator,
+          std::random_access_iterator_tag, clang::AttributeCommonInfo::Kind> {
+  AttrKindIterator() : iterator_adaptor_base(nullptr) {}
+  AttrKindIterator(ParsedAttributesView::const_iterator I)
+      : iterator_adaptor_base(I) {}
+  clang::AttributeCommonInfo::Kind operator*() const { return I->getKind(); }
+};
+
+template <typename AKIterRange>
+CUDAFunctionTarget IdentifyTargetImpl(const AKIterRange &AttrKinds) {
   bool HasHostAttr = false;
   bool HasDeviceAttr = false;
   bool HasGlobalAttr = false;
   bool HasInvalidTargetAttr = false;
-  for (const ParsedAttr &AL : Attrs) {
-    switch (AL.getKind()) {
+  for (const auto &AK : AttrKinds) {
+    switch (AK) {
     case ParsedAttr::AT_CUDAGlobal:
       HasGlobalAttr = true;
       break;
@@ -107,6 +123,18 @@ CUDAFunctionTarget SemaCUDA::IdentifyTarget(const ParsedAttributesView &Attrs) {
   return CUDAFunctionTarget::Host;
 }
 
+} // namespace
+
+CUDAFunctionTarget SemaCUDA::IdentifyTarget(const ParsedAttributesView &Attrs) {
+  return IdentifyTargetImpl(make_range(AttrKindIterator(Attrs.begin()),
+                                       AttrKindIterator(Attrs.end())));
+}
+
+CUDAFunctionTarget SemaCUDA::IdentifyTarget(
+    const SmallVectorImpl<clang::AttributeCommonInfo::Kind> &AttrKinds) {
+  return IdentifyTargetImpl(AttrKinds);
+}
+
 template <typename A>
 static bool hasAttr(const Decl *D, bool IgnoreImplicitAttr) {
   return D->hasAttrs() && llvm::any_of(D->getAttrs(), [&](Attr *Attribute) {
@@ -115,20 +143,65 @@ static bool hasAttr(const Decl *D, bool IgnoreImplicitAttr) {
          });
 }
 
+SemaCUDA::CUDATargetContext::CUDATargetContext(SemaCUDA *S,
+                                               CUDATargetContextKind Kind,
+                                               CUDAFunctionTarget Target)
+    : Kind(Kind), S(S), Target(Target) {}
+
+CUDAFunctionTarget SemaCUDA::CUDATargetContext::getTarget() {
+  TargetQueried = true;
+  return Target;
+}
+
+void SemaCUDA::CUDATargetContext::tryRegisterTargetAttrs(
+    const ParsedAttributesView &Attrs) {
+  if (Kind != CTCK_Declaration)
+    return;
+  for (const auto &A : Attrs) {
+    auto AK = A.getKind();
+    switch (AK) {
+    case ParsedAttr::AT_CUDAGlobal:
+    case ParsedAttr::AT_CUDAHost:
+    case ParsedAttr::AT_CUDADevice:
+    case ParsedAttr::AT_CUDAInvalidTarget:
+      break;
+    default:
+      continue;
+    }
+    AttrKinds.push_back(AK);
+    CUDAFunctionTarget NewTarget = S->IdentifyTarget(AttrKinds);
+    if (TargetQueried && (NewTarget != Target))
+      S->Diag(A.getLoc(), diag::warn_target_specfier_ignored);
+    Target = NewTarget;
+  }
+}
+
 SemaCUDA::CUDATargetContextRAII::CUDATargetContextRAII(
     SemaCUDA &S_, SemaCUDA::CUDATargetContextKind K, Decl *D)
     : S(S_) {
   SavedCtx = S.CurCUDATargetCtx;
-  assert(K == SemaCUDA::CTCK_InitGlobalVar);
-  auto *VD = dyn_cast_or_null<VarDecl>(D);
-  if (VD && VD->hasGlobalStorage() && !VD->isStaticLocal()) {
-    auto Target = CUDAFunctionTarget::Host;
-    if ((hasAttr<CUDADeviceAttr>(VD, /*IgnoreImplicit=*/true) &&
-         !hasAttr<CUDAHostAttr>(VD, /*IgnoreImplicit=*/true)) ||
-        hasAttr<CUDASharedAttr>(VD, /*IgnoreImplicit=*/true) ||
-        hasAttr<CUDAConstantAttr>(VD, /*IgnoreImplicit=*/true))
-      Target = CUDAFunctionTarget::Device;
-    S.CurCUDATargetCtx = {Target, K, VD};
+
+  switch (K) {
+  case SemaCUDA::CTCK_InitGlobalVar: {
+    auto *VD = dyn_cast_or_null<VarDecl>(D);
+    if (VD && VD->hasGlobalStorage() && !VD->isStaticLocal()) {
+      auto Target = CUDAFunctionTarget::Host;
+      if ((hasAttr<CUDADeviceAttr>(VD, /*IgnoreImplicit=*/true) &&
+           !hasAttr<CUDAHostAttr>(VD, /*IgnoreImplicit=*/true)) ||
+          hasAttr<CUDASharedAttr>(VD, /*IgnoreImplicit=*/true) ||
+          hasAttr<CUDAConstantAttr>(VD, /*IgnoreImplicit=*/true))
+        Target = CUDAFunctionTarget::Device;
+      S.CurCUDATargetCtx = CUDATargetContext(&S, K, Target);
+    }
+    break;
+  }
+  case SemaCUDA::CTCK_Declaration:
+    // The target is updated once relevant attributes are parsed. Initialize
+    // with the target used if no attributes are present: Host.
+    S.CurCUDATargetCtx = CUDATargetContext(&S, K, CUDAFunctionTarget::Host);
+    break;
+  default:
+    llvm_unreachable("unexpected context kind");
   }
 }
 
@@ -137,7 +210,7 @@ CUDAFunctionTarget SemaCUDA::IdentifyTarget(const FunctionDecl *D,
                                             bool IgnoreImplicitHDAttr) {
   // Code that lives outside a function gets the target from CurCUDATargetCtx.
   if (D == nullptr)
-    return CurCUDATargetCtx.Target;
+    return CurCUDATargetCtx.getTarget();
 
   if (D->hasAttr<CUDAInvalidTargetAttr>())
     return CUDAFunctionTarget::InvalidTarget;
@@ -232,7 +305,7 @@ SemaCUDA::IdentifyPreference(const FunctionDecl *Caller,
   // trivial ctor/dtor without device attr to be used. Non-trivial ctor/dtor
   // will be diagnosed by checkAllowedInitializer.
   if (Caller == nullptr && CurCUDATargetCtx.Kind == CTCK_InitGlobalVar &&
-      CurCUDATargetCtx.Target == CUDAFunctionTarget::Device &&
+      CurCUDATargetCtx.getTarget() == CUDAFunctionTarget::Device &&
       (isa<CXXConstructorDecl>(Callee) || isa<CXXDestructorDecl>(Callee)))
     return CFP_HostDevice;
 
@@ -297,8 +370,16 @@ SemaCUDA::IdentifyPreference(const FunctionDecl *Caller,
       (CallerTarget == CUDAFunctionTarget::Device &&
        CalleeTarget == CUDAFunctionTarget::Host) ||
       (CallerTarget == CUDAFunctionTarget::Global &&
-       CalleeTarget == CUDAFunctionTarget::Host))
+       CalleeTarget == CUDAFunctionTarget::Host)) {
+    // In declaration contexts outside of function bodies and variable
+    // initializers, tolerate mismatched function targets as long as they are
+    // not codegened.
+    if (CurCUDATargetCtx.Kind == CTCK_Declaration &&
+        !this->SemaRef.getCurFunctionDecl(/*AllowLambda=*/true))
+      return CFP_WrongSide;
+
     return CFP_Never;
+  }
 
   llvm_unreachable("All cases should've been handled by now.");
 }

clang/lib/Sema/SemaOverload.cpp

@@ -10747,7 +10747,7 @@ OverloadCandidateSet::BestViableFunction(Sema &S, SourceLocation Loc,
         llvm::any_of(Candidates, [&](OverloadCandidate *Cand) {
           // Check viable function only.
           return Cand->Viable && Cand->Function &&
-                 S.CUDA().IdentifyPreference(Caller, Cand->Function) ==
+                 S.CUDA().IdentifyPreference(Caller, Cand->Function) >=
                      SemaCUDA::CFP_SameSide;
         });
     if (ContainsSameSideCandidate) {