merge main into amd-staging

Change-Id: I5af07e86eb5e25c35e6277b4385a14fcdb4bc9c8

clang-tools-extra/clang-include-fixer/InMemorySymbolIndex.cpp

@@ -21,7 +21,7 @@ InMemorySymbolIndex::InMemorySymbolIndex(
 
 std::vector<SymbolAndSignals>
 InMemorySymbolIndex::search(llvm::StringRef Identifier) {
-  auto I = LookupTable.find(std::string(Identifier));
+  auto I = LookupTable.find(Identifier);
   if (I != LookupTable.end())
     return I->second;
   return {};

clang-tools-extra/clang-include-fixer/InMemorySymbolIndex.h

@@ -27,7 +27,8 @@ class InMemorySymbolIndex : public SymbolIndex {
   search(llvm::StringRef Identifier) override;
 
 private:
-  std::map<std::string, std::vector<find_all_symbols::SymbolAndSignals>>
+  std::map<std::string, std::vector<find_all_symbols::SymbolAndSignals>,
+           std::less<>>
       LookupTable;
 };
 

clang-tools-extra/clang-tidy/abseil/DurationRewriter.cpp

@@ -244,7 +244,7 @@ std::optional<DurationScale> getScaleForDurationInverse(llvm::StringRef Name) {
        {"ToDoubleNanoseconds", DurationScale::Nanoseconds},
        {"ToInt64Nanoseconds", DurationScale::Nanoseconds}});
 
-  auto ScaleIter = ScaleMap.find(std::string(Name));
+  auto ScaleIter = ScaleMap.find(Name);
   if (ScaleIter == ScaleMap.end())
     return std::nullopt;
 
@@ -260,7 +260,7 @@ std::optional<DurationScale> getScaleForTimeInverse(llvm::StringRef Name) {
        {"ToUnixMicros", DurationScale::Microseconds},
        {"ToUnixNanos", DurationScale::Nanoseconds}});
 
-  auto ScaleIter = ScaleMap.find(std::string(Name));
+  auto ScaleIter = ScaleMap.find(Name);
   if (ScaleIter == ScaleMap.end())
     return std::nullopt;
 

clang/docs/OpenMPSupport.rst

@@ -284,6 +284,8 @@ implementation.
 +------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
 | memory management            | alignment for allocate directive and clause                  | :good:`done`             | D115683                                                               |
 +------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
+| memory management            | 'allocator' modifier for allocate clause                     | :good:`done`             | https://github.com/llvm/llvm-project/pull/114883                      |
++------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
 | memory management            | new memory management routines                               | :none:`unclaimed`        |                                                                       |
 +------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
 | memory management            | changes to omp_alloctrait_key enum                           | :none:`unclaimed`        |                                                                       |

clang/docs/ReleaseNotes.rst

@@ -601,9 +601,6 @@ Bug Fixes to C++ Support
   in certain friend declarations. (#GH93099)
 - Clang now instantiates the correct lambda call operator when a lambda's class type is
   merged across modules. (#GH110401)
-- Clang now uses the correct set of template argument lists when comparing the constraints of
-  out-of-line definitions and member templates explicitly specialized for a given implicit instantiation of
-  a class template. (#GH102320)
 - Fix a crash when parsing a pseudo destructor involving an invalid type. (#GH111460)
 - Fixed an assertion failure when invoking recovery call expressions with explicit attributes
   and undeclared templates. (#GH107047), (#GH49093)

clang/include/clang/AST/DeclTemplate.h

@@ -787,11 +787,15 @@ class RedeclarableTemplateDecl : public TemplateDecl,
                              EntryType *Entry, void *InsertPos);
 
   struct CommonBase {
-    CommonBase() {}
+    CommonBase() : InstantiatedFromMember(nullptr, false) {}
 
     /// The template from which this was most
     /// directly instantiated (or null).
-    RedeclarableTemplateDecl *InstantiatedFromMember = nullptr;
+    ///
+    /// The boolean value indicates whether this template
+    /// was explicitly specialized.
+    llvm::PointerIntPair<RedeclarableTemplateDecl *, 1, bool>
+        InstantiatedFromMember;
 
     /// If non-null, points to an array of specializations (including
     /// partial specializations) known only by their external declaration IDs.
@@ -802,19 +806,14 @@ class RedeclarableTemplateDecl : public TemplateDecl,
   };
 
   /// Pointer to the common data shared by all declarations of this
-  /// template, and a flag indicating if the template is a member
-  /// specialization.
-  mutable llvm::PointerIntPair<CommonBase *, 1, bool> Common;
-
-  CommonBase *getCommonPtrInternal() const { return Common.getPointer(); }
+  /// template.
+  mutable CommonBase *Common = nullptr;
 
   /// Retrieves the "common" pointer shared by all (re-)declarations of
   /// the same template. Calling this routine may implicitly allocate memory
   /// for the common pointer.
   CommonBase *getCommonPtr() const;
 
-  void setCommonPtr(CommonBase *C) const { Common.setPointer(C); }
-
   virtual CommonBase *newCommon(ASTContext &C) const = 0;
 
   // Construct a template decl with name, parameters, and templated element.
@@ -855,12 +854,15 @@ class RedeclarableTemplateDecl : public TemplateDecl,
   /// template<> template<typename T>
   /// struct X<int>::Inner { /* ... */ };
   /// \endcode
-  bool isMemberSpecialization() const { return Common.getInt(); }
+  bool isMemberSpecialization() const {
+    return getCommonPtr()->InstantiatedFromMember.getInt();
+  }
 
   /// Note that this member template is a specialization.
   void setMemberSpecialization() {
-    assert(!isMemberSpecialization() && "already a member specialization");
-    Common.setInt(true);
+    assert(getCommonPtr()->InstantiatedFromMember.getPointer() &&
+           "Only member templates can be member template specializations");
+    getCommonPtr()->InstantiatedFromMember.setInt(true);
   }
 
   /// Retrieve the member template from which this template was
@@ -900,12 +902,12 @@ class RedeclarableTemplateDecl : public TemplateDecl,
   /// void X<T>::f(T, U);
   /// \endcode
   RedeclarableTemplateDecl *getInstantiatedFromMemberTemplate() const {
-    return getCommonPtr()->InstantiatedFromMember;
+    return getCommonPtr()->InstantiatedFromMember.getPointer();
   }
 
   void setInstantiatedFromMemberTemplate(RedeclarableTemplateDecl *TD) {
-    assert(!getCommonPtr()->InstantiatedFromMember);
-    getCommonPtr()->InstantiatedFromMember = TD;
+    assert(!getCommonPtr()->InstantiatedFromMember.getPointer());
+    getCommonPtr()->InstantiatedFromMember.setPointer(TD);
   }
 
   /// Retrieve the "injected" template arguments that correspond to the
@@ -1955,7 +1957,13 @@ class ClassTemplateSpecializationDecl : public CXXRecordDecl,
   /// specialization which was specialized by this.
   llvm::PointerUnion<ClassTemplateDecl *,
                      ClassTemplatePartialSpecializationDecl *>
-  getSpecializedTemplateOrPartial() const;
+  getSpecializedTemplateOrPartial() const {
+    if (const auto *PartialSpec =
+            SpecializedTemplate.dyn_cast<SpecializedPartialSpecialization *>())
+      return PartialSpec->PartialSpecialization;
+
+    return SpecializedTemplate.get<ClassTemplateDecl*>();
+  }
 
   /// Retrieve the set of template arguments that should be used
   /// to instantiate members of the class template or class template partial
@@ -1981,8 +1989,6 @@ class ClassTemplateSpecializationDecl : public CXXRecordDecl,
   /// template arguments have been deduced.
   void setInstantiationOf(ClassTemplatePartialSpecializationDecl *PartialSpec,
                           const TemplateArgumentList *TemplateArgs) {
-    assert(!isa<ClassTemplatePartialSpecializationDecl>(this) &&
-           "A partial specialization cannot be instantiated from a template");
     assert(!SpecializedTemplate.is<SpecializedPartialSpecialization*>() &&
            "Already set to a class template partial specialization!");
     auto *PS = new (getASTContext()) SpecializedPartialSpecialization();
@@ -1994,8 +2000,6 @@ class ClassTemplateSpecializationDecl : public CXXRecordDecl,
   /// Note that this class template specialization is an instantiation
   /// of the given class template.
   void setInstantiationOf(ClassTemplateDecl *TemplDecl) {
-    assert(!isa<ClassTemplatePartialSpecializationDecl>(this) &&
-           "A partial specialization cannot be instantiated from a template");
     assert(!SpecializedTemplate.is<SpecializedPartialSpecialization*>() &&
            "Previously set to a class template partial specialization!");
     SpecializedTemplate = TemplDecl;
@@ -2189,11 +2193,18 @@ class ClassTemplatePartialSpecializationDecl
   /// struct X<int>::Inner<T*> { /* ... */ };
   /// \endcode
   bool isMemberSpecialization() const {
-    return InstantiatedFromMember.getInt();
+    const auto *First =
+        cast<ClassTemplatePartialSpecializationDecl>(getFirstDecl());
+    return First->InstantiatedFromMember.getInt();
   }
 
   /// Note that this member template is a specialization.
-  void setMemberSpecialization() { return InstantiatedFromMember.setInt(true); }
+  void setMemberSpecialization() {
+    auto *First = cast<ClassTemplatePartialSpecializationDecl>(getFirstDecl());
+    assert(First->InstantiatedFromMember.getPointer() &&
+           "Only member templates can be member template specializations");
+    return First->InstantiatedFromMember.setInt(true);
+  }
 
   /// Retrieves the injected specialization type for this partial
   /// specialization.  This is not the same as the type-decl-type for
@@ -2263,6 +2274,8 @@ class ClassTemplateDecl : public RedeclarableTemplateDecl {
     return static_cast<Common *>(RedeclarableTemplateDecl::getCommonPtr());
   }
 
+  void setCommonPtr(Common *C) { RedeclarableTemplateDecl::Common = C; }
+
 public:
 
   friend class ASTDeclReader;
@@ -2713,7 +2726,13 @@ class VarTemplateSpecializationDecl : public VarDecl,
   /// Retrieve the variable template or variable template partial
   /// specialization which was specialized by this.
   llvm::PointerUnion<VarTemplateDecl *, VarTemplatePartialSpecializationDecl *>
-  getSpecializedTemplateOrPartial() const;
+  getSpecializedTemplateOrPartial() const {
+    if (const auto *PartialSpec =
+            SpecializedTemplate.dyn_cast<SpecializedPartialSpecialization *>())
+      return PartialSpec->PartialSpecialization;
+
+    return SpecializedTemplate.get<VarTemplateDecl *>();
+  }
 
   /// Retrieve the set of template arguments that should be used
   /// to instantiate the initializer of the variable template or variable
@@ -2739,8 +2758,6 @@ class VarTemplateSpecializationDecl : public VarDecl,
   /// template arguments have been deduced.
   void setInstantiationOf(VarTemplatePartialSpecializationDecl *PartialSpec,
                           const TemplateArgumentList *TemplateArgs) {
-    assert(!isa<VarTemplatePartialSpecializationDecl>(this) &&
-           "A partial specialization cannot be instantiated from a template");
     assert(!SpecializedTemplate.is<SpecializedPartialSpecialization *>() &&
            "Already set to a variable template partial specialization!");
     auto *PS = new (getASTContext()) SpecializedPartialSpecialization();
@@ -2752,8 +2769,6 @@ class VarTemplateSpecializationDecl : public VarDecl,
   /// Note that this variable template specialization is an instantiation
   /// of the given variable template.
   void setInstantiationOf(VarTemplateDecl *TemplDecl) {
-    assert(!isa<VarTemplatePartialSpecializationDecl>(this) &&
-           "A partial specialization cannot be instantiated from a template");
     assert(!SpecializedTemplate.is<SpecializedPartialSpecialization *>() &&
            "Previously set to a variable template partial specialization!");
     SpecializedTemplate = TemplDecl;
@@ -2944,11 +2959,18 @@ class VarTemplatePartialSpecializationDecl
   /// U* X<int>::Inner<T*> = (T*)(0) + 1;
   /// \endcode
   bool isMemberSpecialization() const {
-    return InstantiatedFromMember.getInt();
+    const auto *First =
+        cast<VarTemplatePartialSpecializationDecl>(getFirstDecl());
+    return First->InstantiatedFromMember.getInt();
   }
 
   /// Note that this member template is a specialization.
-  void setMemberSpecialization() { return InstantiatedFromMember.setInt(true); }
+  void setMemberSpecialization() {
+    auto *First = cast<VarTemplatePartialSpecializationDecl>(getFirstDecl());
+    assert(First->InstantiatedFromMember.getPointer() &&
+           "Only member templates can be member template specializations");
+    return First->InstantiatedFromMember.setInt(true);
+  }
 
   SourceRange getSourceRange() const override LLVM_READONLY;
 
@@ -3119,9 +3141,6 @@ class VarTemplateDecl : public RedeclarableTemplateDecl {
     return makeSpecIterator(getSpecializations(), true);
   }
 
-  /// Merge \p Prev with our RedeclarableTemplateDecl::Common.
-  void mergePrevDecl(VarTemplateDecl *Prev);
-
   // Implement isa/cast/dyncast support
   static bool classof(const Decl *D) { return classofKind(D->getKind()); }
   static bool classofKind(Kind K) { return K == VarTemplate; }

clang/include/clang/AST/SYCLKernelInfo.h

@@ -13,6 +13,7 @@
 #ifndef LLVM_CLANG_AST_SYCLKERNELINFO_H
 #define LLVM_CLANG_AST_SYCLKERNELINFO_H
 
+#include "clang/AST/CanonicalType.h"
 #include "clang/AST/Decl.h"
 #include "clang/AST/Type.h"
 

clang/include/clang/Basic/Builtins.td

@@ -4798,6 +4798,12 @@ def HLSLDot4AddI8Packed : LangBuiltin<"HLSL_LANG"> {
   let Prototype = "int(unsigned int, unsigned int, int)";
 }
 
+def HLSLFirstBitHigh : LangBuiltin<"HLSL_LANG"> {
+  let Spellings = ["__builtin_hlsl_elementwise_firstbithigh"];
+  let Attributes = [NoThrow, Const];
+  let Prototype = "void(...)";
+}
+
 def HLSLFrac : LangBuiltin<"HLSL_LANG"> {
   let Spellings = ["__builtin_hlsl_elementwise_frac"];
   let Attributes = [NoThrow, Const];

clang/include/clang/Basic/BuiltinsAMDGPU.def

@@ -439,15 +439,15 @@ TARGET_BUILTIN(__builtin_amdgcn_s_sleep_var, "vUi", "n", "gfx12-insts")
 TARGET_BUILTIN(__builtin_amdgcn_permlane16_var,  "UiUiUiUiIbIb", "nc", "gfx12-insts")
 TARGET_BUILTIN(__builtin_amdgcn_permlanex16_var, "UiUiUiUiIbIb", "nc", "gfx12-insts")
 TARGET_BUILTIN(__builtin_amdgcn_s_barrier_signal, "vIi", "n", "gfx12-insts")
-TARGET_BUILTIN(__builtin_amdgcn_s_barrier_signal_var, "vi", "n", "gfx12-insts")
+TARGET_BUILTIN(__builtin_amdgcn_s_barrier_signal_var, "vv*i", "n", "gfx12-insts")
 TARGET_BUILTIN(__builtin_amdgcn_s_barrier_wait, "vIs", "n", "gfx12-insts")
 TARGET_BUILTIN(__builtin_amdgcn_s_barrier_signal_isfirst, "bIi", "n", "gfx12-insts")
-TARGET_BUILTIN(__builtin_amdgcn_s_barrier_signal_isfirst_var, "bi", "n", "gfx12-insts")
-TARGET_BUILTIN(__builtin_amdgcn_s_barrier_init, "vii", "n", "gfx12-insts")
-TARGET_BUILTIN(__builtin_amdgcn_s_barrier_join, "vi", "n", "gfx12-insts")
-TARGET_BUILTIN(__builtin_amdgcn_s_wakeup_barrier, "vi", "n", "gfx12-insts")
-TARGET_BUILTIN(__builtin_amdgcn_s_barrier_leave, "b", "n", "gfx12-insts")
+TARGET_BUILTIN(__builtin_amdgcn_s_barrier_init, "vv*i", "n", "gfx12-insts")
+TARGET_BUILTIN(__builtin_amdgcn_s_barrier_join, "vv*", "n", "gfx12-insts")
+TARGET_BUILTIN(__builtin_amdgcn_s_wakeup_barrier, "vv*", "n", "gfx12-insts")
+TARGET_BUILTIN(__builtin_amdgcn_s_barrier_leave, "vIs", "n", "gfx12-insts")
 TARGET_BUILTIN(__builtin_amdgcn_s_get_barrier_state, "Uii", "n", "gfx12-insts")
+TARGET_BUILTIN(__builtin_amdgcn_s_get_named_barrier_state, "Uiv*", "n", "gfx12-insts")
 TARGET_BUILTIN(__builtin_amdgcn_s_prefetch_data, "vvC*Ui", "nc", "gfx12-insts")
 TARGET_BUILTIN(__builtin_amdgcn_s_buffer_prefetch_data, "vQbIiUi", "nc", "gfx12-insts")
 

clang/include/clang/Sema/Sema.h

@@ -11328,9 +11328,9 @@ class Sema final : public SemaBase {
       CXXScopeSpec &SS, IdentifierInfo *Name, SourceLocation NameLoc,
       const ParsedAttributesView &Attr, TemplateParameterList *TemplateParams,
       AccessSpecifier AS, SourceLocation ModulePrivateLoc,
-      SourceLocation FriendLoc,
-      ArrayRef<TemplateParameterList *> OuterTemplateParamLists,
-      bool IsMemberSpecialization, SkipBodyInfo *SkipBody = nullptr);
+      SourceLocation FriendLoc, unsigned NumOuterTemplateParamLists,
+      TemplateParameterList **OuterTemplateParamLists,
+      SkipBodyInfo *SkipBody = nullptr);
 
   /// Translates template arguments as provided by the parser
   /// into template arguments used by semantic analysis.
@@ -11369,8 +11369,7 @@ class Sema final : public SemaBase {
   DeclResult ActOnVarTemplateSpecialization(
       Scope *S, Declarator &D, TypeSourceInfo *DI, LookupResult &Previous,
       SourceLocation TemplateKWLoc, TemplateParameterList *TemplateParams,
-      StorageClass SC, bool IsPartialSpecialization,
-      bool IsMemberSpecialization);
+      StorageClass SC, bool IsPartialSpecialization);
 
   /// Get the specialization of the given variable template corresponding to
   /// the specified argument list, or a null-but-valid result if the arguments
@@ -13012,14 +13011,28 @@ class Sema final : public SemaBase {
   /// dealing with a specialization. This is only relevant for function
   /// template specializations.
   ///
+  /// \param Pattern If non-NULL, indicates the pattern from which we will be
+  /// instantiating the definition of the given declaration, \p ND. This is
+  /// used to determine the proper set of template instantiation arguments for
+  /// friend function template specializations.
+  ///
   /// \param ForConstraintInstantiation when collecting arguments,
   /// ForConstraintInstantiation indicates we should continue looking when
   /// encountering a lambda generic call operator, and continue looking for
   /// arguments on an enclosing class template.
+  ///
+  /// \param SkipForSpecialization when specified, any template specializations
+  /// in a traversal would be ignored.
+  /// \param ForDefaultArgumentSubstitution indicates we should continue looking
+  /// when encountering a specialized member function template, rather than
+  /// returning immediately.
   MultiLevelTemplateArgumentList getTemplateInstantiationArgs(
       const NamedDecl *D, const DeclContext *DC = nullptr, bool Final = false,
       std::optional<ArrayRef<TemplateArgument>> Innermost = std::nullopt,
-      bool RelativeToPrimary = false, bool ForConstraintInstantiation = false);
+      bool RelativeToPrimary = false, const FunctionDecl *Pattern = nullptr,
+      bool ForConstraintInstantiation = false,
+      bool SkipForSpecialization = false,
+      bool ForDefaultArgumentSubstitution = false);
 
   /// RAII object to handle the state changes required to synthesize
   /// a function body.

clang/lib/AST/ASTImporter.cpp

@@ -6190,8 +6190,7 @@ ExpectedDecl ASTNodeImporter::VisitClassTemplateDecl(ClassTemplateDecl *D) {
 ExpectedDecl ASTNodeImporter::VisitClassTemplateSpecializationDecl(
                                           ClassTemplateSpecializationDecl *D) {
   ClassTemplateDecl *ClassTemplate;
-  if (Error Err = importInto(ClassTemplate,
-                             D->getSpecializedTemplate()->getCanonicalDecl()))
+  if (Error Err = importInto(ClassTemplate, D->getSpecializedTemplate()))
     return std::move(Err);
 
   // Import the context of this declaration.