Vendor import of llvm-project main llvmorg-18-init-15088-gd14ee76181fb.

clang/include/clang-c/BuildSystem.h

@@ -95,7 +95,7 @@ CINDEX_LINKAGE void clang_free(void *buffer);
 CINDEX_LINKAGE void clang_VirtualFileOverlay_dispose(CXVirtualFileOverlay);
 
 /**
- * Object encapsulating information about a module.map file.
+ * Object encapsulating information about a module.modulemap file.
  */
 typedef struct CXModuleMapDescriptorImpl *CXModuleMapDescriptor;
 
@@ -109,7 +109,7 @@ CINDEX_LINKAGE CXModuleMapDescriptor
 clang_ModuleMapDescriptor_create(unsigned options);
 
 /**
- * Sets the framework module name that the module.map describes.
+ * Sets the framework module name that the module.modulemap describes.
  * \returns 0 for success, non-zero to indicate an error.
  */
 CINDEX_LINKAGE enum CXErrorCode
@@ -117,7 +117,7 @@ clang_ModuleMapDescriptor_setFrameworkModuleName(CXModuleMapDescriptor,
                                                  const char *name);
 
 /**
- * Sets the umbrella header name that the module.map describes.
+ * Sets the umbrella header name that the module.modulemap describes.
  * \returns 0 for success, non-zero to indicate an error.
  */
 CINDEX_LINKAGE enum CXErrorCode

clang/include/clang/AST/Type.h

@@ -36,6 +36,7 @@
 #include "llvm/ADT/FoldingSet.h"
 #include "llvm/ADT/PointerIntPair.h"
 #include "llvm/ADT/PointerUnion.h"
+#include "llvm/ADT/STLForwardCompat.h"
 #include "llvm/ADT/StringRef.h"
 #include "llvm/ADT/Twine.h"
 #include "llvm/ADT/iterator_range.h"
@@ -7514,7 +7515,7 @@ inline const Type *Type::getPointeeOrArrayElementType() const {
 /// spaces into a diagnostic with <<.
 inline const StreamingDiagnostic &operator<<(const StreamingDiagnostic &PD,
                                              LangAS AS) {
-  PD.AddTaggedVal(static_cast<std::underlying_type_t<LangAS>>(AS),
+  PD.AddTaggedVal(llvm::to_underlying(AS),
                   DiagnosticsEngine::ArgumentKind::ak_addrspace);
   return PD;
 }

clang/include/clang/Analysis/Analyses/UnsafeBufferUsageGadgets.def

@@ -36,6 +36,7 @@ FIXABLE_GADGET(PointerDereference)
 FIXABLE_GADGET(UPCAddressofArraySubscript) // '&DRE[any]' in an Unspecified Pointer Context
 FIXABLE_GADGET(UPCStandalonePointer)
 FIXABLE_GADGET(UPCPreIncrement)            // '++Ptr' in an Unspecified Pointer Context
+FIXABLE_GADGET(UUCAddAssign)            // 'Ptr += n' in an Unspecified Untyped Context
 FIXABLE_GADGET(PointerAssignment)
 FIXABLE_GADGET(PointerInit)
 

clang/include/clang/Basic/Attr.td

@@ -455,6 +455,9 @@ def TargetMicrosoftCXXABI : TargetArch<["x86", "x86_64", "arm", "thumb", "aarch6
 def TargetELF : TargetSpec {
   let ObjectFormats = ["ELF"];
 }
+def TargetELFOrMachO : TargetSpec {
+  let ObjectFormats = ["ELF", "MachO"];
+}
 
 def TargetSupportsInitPriority : TargetSpec {
   let CustomCode = [{ !Target.getTriple().isOSzOS() }];
@@ -1665,7 +1668,7 @@ def IBOutletCollection : InheritableAttr {
   let Documentation = [Undocumented];
 }
 
-def IFunc : Attr, TargetSpecificAttr<TargetELF> {
+def IFunc : Attr, TargetSpecificAttr<TargetELFOrMachO> {
   let Spellings = [GCC<"ifunc">];
   let Args = [StringArgument<"Resolver">];
   let Subjects = SubjectList<[Function]>;
@@ -2212,7 +2215,8 @@ def NotTailCalled : InheritableAttr {
 def : MutualExclusions<[AlwaysInline, NotTailCalled]>;
 
 def NoStackProtector : InheritableAttr {
-  let Spellings = [Clang<"no_stack_protector">, Declspec<"safebuffers">];
+  let Spellings = [Clang<"no_stack_protector">, CXX11<"gnu", "no_stack_protector">,
+                   C23<"gnu", "no_stack_protector">, Declspec<"safebuffers">];
   let Subjects = SubjectList<[Function]>;
   let Documentation = [NoStackProtectorDocs];
   let SimpleHandler = 1;
@@ -2792,9 +2796,10 @@ def SwiftAsyncError : InheritableAttr {
   let Documentation = [SwiftAsyncErrorDocs];
 }
 
-def Suppress : StmtAttr {
-  let Spellings = [CXX11<"gsl", "suppress">];
+def Suppress : DeclOrStmtAttr {
+  let Spellings = [CXX11<"gsl", "suppress">, Clang<"suppress">];
   let Args = [VariadicStringArgument<"DiagnosticIdentifiers">];
+  let Accessors = [Accessor<"isGSL", [CXX11<"gsl", "suppress">]>];
   let Documentation = [SuppressDocs];
 }
 
@@ -2878,7 +2883,7 @@ def Target : InheritableAttr {
 
       for (auto &Feature : AttrFeatures) {
         Feature = Feature.trim();
-        if (Feature.startswith("arch="))
+        if (Feature.starts_with("arch="))
           return Feature.drop_front(sizeof("arch=") - 1);
       }
       return "";
@@ -2896,8 +2901,8 @@ def Target : InheritableAttr {
       for (auto &Feature : AttrFeatures) {
         Feature = Feature.trim();
 
-        if (!Feature.startswith("no-") && !Feature.startswith("arch=") &&
-            !Feature.startswith("fpmath=") && !Feature.startswith("tune="))
+        if (!Feature.starts_with("no-") && !Feature.starts_with("arch=") &&
+            !Feature.starts_with("fpmath=") && !Feature.starts_with("tune="))
           Out.push_back(Feature);
       }
     }
@@ -4258,7 +4263,8 @@ def HLSLResource : InheritableAttr {
                             "StructuredBuffer", "CBuffer", "Sampler",
                             "TBuffer", "RTAccelerationStructure",
                             "FeedbackTexture2D", "FeedbackTexture2DArray"],
-                           /*opt=*/0, /*fake=*/0, /*isExternalType=*/1>
+                           /*opt=*/0, /*fake=*/0, /*isExternalType=*/1>,
+              DefaultBoolArgument<"isROV", /*default=*/0>
               ];
   let Documentation = [InternalOnly];
 }

clang/include/clang/Basic/AttrDocs.td

@@ -5267,7 +5267,74 @@ the ``int`` parameter is the one that represents the error.
 def SuppressDocs : Documentation {
   let Category = DocCatStmt;
   let Content = [{
-The ``[[gsl::suppress]]`` attribute suppresses specific
+The ``suppress`` attribute suppresses unwanted warnings coming from static
+analysis tools such as the Clang Static Analyzer. The tool will not report
+any issues in source code annotated with the attribute.
+
+The attribute cannot be used to suppress traditional Clang warnings, because
+many such warnings are emitted before the attribute is fully parsed.
+Consider using ``#pragma clang diagnostic`` to control such diagnostics,
+as described in `Controlling Diagnostics via Pragmas
+<https://clang.llvm.org/docs/UsersManual.html#controlling-diagnostics-via-pragmas>`_.
+
+The ``suppress`` attribute can be placed on an individual statement in order to
+suppress warnings about undesirable behavior occurring at that statement:
+
+.. code-block:: c++
+
+  int foo() {
+    int *x = nullptr;
+    ...
+    [[clang::suppress]]
+    return *x;  // null pointer dereference warning suppressed here
+  }
+
+Putting the attribute on a compound statement suppresses all warnings in scope:
+
+.. code-block:: c++
+
+  int foo() {
+    [[clang::suppress]] {
+      int *x = nullptr;
+      ...
+      return *x;  // warnings suppressed in the entire scope
+    }
+  }
+
+Some static analysis warnings are accompanied by one or more notes, and the
+line of code against which the warning is emitted isn't necessarily the best
+for suppression purposes. In such cases the tools are allowed to implement
+additional ways to suppress specific warnings based on the attribute attached
+to a note location.
+
+For example, the Clang Static Analyzer suppresses memory leak warnings when
+the suppression attribute is placed at the allocation site (highlited by
+a "note: memory is allocated"), which may be different from the line of code
+at which the program "loses track" of the pointer (where the warning
+is ultimately emitted):
+
+.. code-block:: c
+
+  int bar1(bool coin_flip) {
+    __attribute__((suppress))
+    int *result = (int *)malloc(sizeof(int));
+    if (coin_flip)
+      return 1;  // warning about this leak path is suppressed
+
+    return *result;  // warning about this leak path is also suppressed
+  }
+
+  int bar2(bool coin_flip) {
+    int *result = (int *)malloc(sizeof(int));
+    if (coin_flip)
+      return 1;  // leak warning on this path NOT suppressed
+
+    __attribute__((suppress))
+    return *result;  // leak warning is suppressed only on this path
+  }
+
+
+When written as ``[[gsl::suppress]]``, this attribute suppresses specific
 clang-tidy diagnostics for rules of the `C++ Core Guidelines`_ in a portable
 way. The attribute can be attached to declarations, statements, and at
 namespace scope.
@@ -5468,7 +5535,9 @@ considered inline.
 Not all targets support this attribute. ELF target support depends on both the
 linker and runtime linker, and is available in at least lld 4.0 and later,
 binutils 2.20.1 and later, glibc v2.11.1 and later, and FreeBSD 9.1 and later.
-Non-ELF targets currently do not support this attribute.
+Mach-O targets support it, but with slightly different semantics: the resolver
+is run at first call, instead of at load time by the runtime linker. Targets
+other than ELF and Mach-O currently do not support this attribute.
   }];
 }
 

clang/include/clang/Basic/BuiltinsAMDGPU.def

@@ -406,5 +406,23 @@ TARGET_BUILTIN(__builtin_amdgcn_cvt_pk_fp8_f32, "iffiIb", "nc", "fp8-insts")
 TARGET_BUILTIN(__builtin_amdgcn_cvt_sr_bf8_f32, "ifiiIi", "nc", "fp8-insts")
 TARGET_BUILTIN(__builtin_amdgcn_cvt_sr_fp8_f32, "ifiiIi", "nc", "fp8-insts")
 
+//===----------------------------------------------------------------------===//
+// GFX12+ only builtins.
+//===----------------------------------------------------------------------===//
+
+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_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_get_barrier_state, "Uii", "n", "gfx12-insts")
+
+
 #undef BUILTIN
 #undef TARGET_BUILTIN

clang/include/clang/Basic/BuiltinsNVPTX.def

@@ -26,7 +26,9 @@
 #pragma push_macro("SM_87")
 #pragma push_macro("SM_89")
 #pragma push_macro("SM_90")
-#define SM_90 "sm_90"
+#pragma push_macro("SM_90a")
+#define SM_90a "sm_90a"
+#define SM_90 "sm_90|" SM_90a
 #define SM_89 "sm_89|" SM_90
 #define SM_87 "sm_87|" SM_89
 #define SM_86 "sm_86|" SM_87
@@ -56,7 +58,11 @@
 #pragma push_macro("PTX78")
 #pragma push_macro("PTX80")
 #pragma push_macro("PTX81")
-#define PTX81 "ptx81"
+#pragma push_macro("PTX82")
+#pragma push_macro("PTX83")
+#define PTX83 "ptx83"
+#define PTX82 "ptx82|" PTX83
+#define PTX81 "ptx81|" PTX82
 #define PTX80 "ptx80|" PTX81
 #define PTX78 "ptx78|" PTX80
 #define PTX77 "ptx77|" PTX78
@@ -1055,6 +1061,7 @@ TARGET_BUILTIN(__nvvm_getctarank_shared_cluster, "iv*3", "", AND(SM_90,PTX78))
 #pragma pop_macro("SM_87")
 #pragma pop_macro("SM_89")
 #pragma pop_macro("SM_90")
+#pragma pop_macro("SM_90a")
 #pragma pop_macro("PTX42")
 #pragma pop_macro("PTX60")
 #pragma pop_macro("PTX61")
@@ -1072,3 +1079,5 @@ TARGET_BUILTIN(__nvvm_getctarank_shared_cluster, "iv*3", "", AND(SM_90,PTX78))
 #pragma pop_macro("PTX78")
 #pragma pop_macro("PTX80")
 #pragma pop_macro("PTX81")
+#pragma pop_macro("PTX82")
+#pragma pop_macro("PTX83")

clang/include/clang/Basic/BuiltinsSystemZ.def

@@ -64,14 +64,14 @@ TARGET_BUILTIN(__builtin_s390_vupllh, "V4UiV8Us", "nc", "vector")
 TARGET_BUILTIN(__builtin_s390_vupllf, "V2ULLiV4Ui", "nc", "vector")
 
 // Vector integer instructions (chapter 22 of the PoP)
-TARGET_BUILTIN(__builtin_s390_vaq, "V16UcV16UcV16Uc", "nc", "vector")
-TARGET_BUILTIN(__builtin_s390_vacq, "V16UcV16UcV16UcV16Uc", "nc", "vector")
+TARGET_BUILTIN(__builtin_s390_vaq, "SLLLiSLLLiSLLLi", "nc", "vector")
+TARGET_BUILTIN(__builtin_s390_vacq, "ULLLiULLLiULLLiULLLi", "nc", "vector")
 TARGET_BUILTIN(__builtin_s390_vaccb, "V16UcV16UcV16Uc", "nc", "vector")
 TARGET_BUILTIN(__builtin_s390_vacch, "V8UsV8UsV8Us", "nc", "vector")
 TARGET_BUILTIN(__builtin_s390_vaccf, "V4UiV4UiV4Ui", "nc", "vector")
 TARGET_BUILTIN(__builtin_s390_vaccg, "V2ULLiV2ULLiV2ULLi", "nc", "vector")
-TARGET_BUILTIN(__builtin_s390_vaccq, "V16UcV16UcV16Uc", "nc", "vector")
-TARGET_BUILTIN(__builtin_s390_vacccq, "V16UcV16UcV16UcV16Uc", "nc", "vector")
+TARGET_BUILTIN(__builtin_s390_vaccq, "ULLLiULLLiULLLi", "nc", "vector")
+TARGET_BUILTIN(__builtin_s390_vacccq, "ULLLiULLLiULLLiULLLi", "nc", "vector")
 TARGET_BUILTIN(__builtin_s390_vavgb, "V16ScV16ScV16Sc", "nc", "vector")
 TARGET_BUILTIN(__builtin_s390_vavgh, "V8SsV8SsV8Ss", "nc", "vector")
 TARGET_BUILTIN(__builtin_s390_vavgf, "V4SiV4SiV4Si", "nc", "vector")
@@ -116,11 +116,11 @@ TARGET_BUILTIN(__builtin_s390_verllvg, "V2ULLiV2ULLiV2ULLi", "nc", "vector")
 TARGET_BUILTIN(__builtin_s390_vgfmb, "V8UsV16UcV16Uc", "nc", "vector")
 TARGET_BUILTIN(__builtin_s390_vgfmh, "V4UiV8UsV8Us", "nc", "vector")
 TARGET_BUILTIN(__builtin_s390_vgfmf, "V2ULLiV4UiV4Ui", "nc", "vector")
-TARGET_BUILTIN(__builtin_s390_vgfmg, "V16UcV2ULLiV2ULLi", "nc", "vector")
+TARGET_BUILTIN(__builtin_s390_vgfmg, "ULLLiV2ULLiV2ULLi", "nc", "vector")
 TARGET_BUILTIN(__builtin_s390_vgfmab, "V8UsV16UcV16UcV8Us", "nc", "vector")
 TARGET_BUILTIN(__builtin_s390_vgfmah, "V4UiV8UsV8UsV4Ui", "nc", "vector")
 TARGET_BUILTIN(__builtin_s390_vgfmaf, "V2ULLiV4UiV4UiV2ULLi", "nc", "vector")
-TARGET_BUILTIN(__builtin_s390_vgfmag, "V16UcV2ULLiV2ULLiV16Uc", "nc", "vector")
+TARGET_BUILTIN(__builtin_s390_vgfmag, "ULLLiV2ULLiV2ULLiULLLi", "nc", "vector")
 TARGET_BUILTIN(__builtin_s390_vmahb, "V16ScV16ScV16ScV16Sc", "nc", "vector")
 TARGET_BUILTIN(__builtin_s390_vmahh, "V8SsV8SsV8SsV8Ss", "nc", "vector")
 TARGET_BUILTIN(__builtin_s390_vmahf, "V4SiV4SiV4SiV4Si", "nc", "vector")
@@ -161,14 +161,14 @@ TARGET_BUILTIN(__builtin_s390_vpopctb, "V16UcV16Uc", "nc", "vector")
 TARGET_BUILTIN(__builtin_s390_vpopcth, "V8UsV8Us", "nc", "vector")
 TARGET_BUILTIN(__builtin_s390_vpopctf, "V4UiV4Ui", "nc", "vector")
 TARGET_BUILTIN(__builtin_s390_vpopctg, "V2ULLiV2ULLi", "nc", "vector")
-TARGET_BUILTIN(__builtin_s390_vsq, "V16UcV16UcV16Uc", "nc", "vector")
-TARGET_BUILTIN(__builtin_s390_vsbcbiq, "V16UcV16UcV16UcV16Uc", "nc", "vector")
-TARGET_BUILTIN(__builtin_s390_vsbiq, "V16UcV16UcV16UcV16Uc", "nc", "vector")
+TARGET_BUILTIN(__builtin_s390_vsq, "SLLLiSLLLiSLLLi", "nc", "vector")
+TARGET_BUILTIN(__builtin_s390_vsbcbiq, "ULLLiULLLiULLLiULLLi", "nc", "vector")
+TARGET_BUILTIN(__builtin_s390_vsbiq, "ULLLiULLLiULLLiULLLi", "nc", "vector")
 TARGET_BUILTIN(__builtin_s390_vscbib, "V16UcV16UcV16Uc", "nc", "vector")
 TARGET_BUILTIN(__builtin_s390_vscbih, "V8UsV8UsV8Us", "nc", "vector")
 TARGET_BUILTIN(__builtin_s390_vscbif, "V4UiV4UiV4Ui", "nc", "vector")
 TARGET_BUILTIN(__builtin_s390_vscbig, "V2ULLiV2ULLiV2ULLi", "nc", "vector")
-TARGET_BUILTIN(__builtin_s390_vscbiq, "V16UcV16UcV16Uc", "nc", "vector")
+TARGET_BUILTIN(__builtin_s390_vscbiq, "ULLLiULLLiULLLi", "nc", "vector")
 TARGET_BUILTIN(__builtin_s390_vsl, "V16UcV16UcV16Uc", "nc", "vector")
 TARGET_BUILTIN(__builtin_s390_vslb, "V16UcV16UcV16Uc", "nc", "vector")
 TARGET_BUILTIN(__builtin_s390_vsldb, "V16UcV16UcV16UcIi", "nc", "vector")
@@ -180,8 +180,8 @@ TARGET_BUILTIN(__builtin_s390_vsumb, "V4UiV16UcV16Uc", "nc", "vector")
 TARGET_BUILTIN(__builtin_s390_vsumh, "V4UiV8UsV8Us", "nc", "vector")
 TARGET_BUILTIN(__builtin_s390_vsumgh, "V2ULLiV8UsV8Us", "nc", "vector")
 TARGET_BUILTIN(__builtin_s390_vsumgf, "V2ULLiV4UiV4Ui", "nc", "vector")
-TARGET_BUILTIN(__builtin_s390_vsumqf, "V16UcV4UiV4Ui", "nc", "vector")
-TARGET_BUILTIN(__builtin_s390_vsumqg, "V16UcV2ULLiV2ULLi", "nc", "vector")
+TARGET_BUILTIN(__builtin_s390_vsumqf, "ULLLiV4UiV4Ui", "nc", "vector")
+TARGET_BUILTIN(__builtin_s390_vsumqg, "ULLLiV2ULLiV2ULLi", "nc", "vector")
 TARGET_BUILTIN(__builtin_s390_vtm, "iV16UcV16Uc", "nc", "vector")
 
 // Vector string instructions (chapter 23 of the PoP)
@@ -256,7 +256,7 @@ TARGET_BUILTIN(__builtin_s390_vftcidb, "V2SLLiV2dIii*", "nc", "vector")
 TARGET_BUILTIN(__builtin_s390_vlrlr, "V16ScUivC*", "", "vector-enhancements-1")
 TARGET_BUILTIN(__builtin_s390_vstrlr, "vV16ScUiv*", "", "vector-enhancements-1")
 TARGET_BUILTIN(__builtin_s390_vbperm, "V2ULLiV16UcV16Uc", "nc", "vector-enhancements-1")
-TARGET_BUILTIN(__builtin_s390_vmslg, "V16UcV2ULLiV2ULLiV16UcIi", "nc", "vector-enhancements-1")
+TARGET_BUILTIN(__builtin_s390_vmslg, "ULLLiV2ULLiV2ULLiULLLiIi", "nc", "vector-enhancements-1")
 TARGET_BUILTIN(__builtin_s390_vfmaxdb, "V2dV2dV2dIi", "nc", "vector-enhancements-1")
 TARGET_BUILTIN(__builtin_s390_vfmindb, "V2dV2dV2dIi", "nc", "vector-enhancements-1")
 TARGET_BUILTIN(__builtin_s390_vfnmadb, "V2dV2dV2dV2d", "nc", "vector-enhancements-1")

clang/include/clang/Basic/BuiltinsX86.def

@@ -979,7 +979,7 @@ TARGET_BUILTIN(__builtin_ia32_scatterpfqps, "vUcV8Oiv*IiIi", "nV:512:", "avx512p
 TARGET_BUILTIN(__builtin_ia32_knotqi, "UcUc", "nc", "avx512dq")
 TARGET_BUILTIN(__builtin_ia32_knothi, "UsUs", "nc", "avx512f")
 TARGET_BUILTIN(__builtin_ia32_knotsi, "UiUi", "nc", "avx512bw")
-TARGET_BUILTIN(__builtin_ia32_knotdi, "UOiUOi", "nc", "avx512bw,evex512")
+TARGET_BUILTIN(__builtin_ia32_knotdi, "UOiUOi", "nc", "avx512bw")
 
 TARGET_BUILTIN(__builtin_ia32_cmpb128_mask, "UsV16cV16cIiUs", "ncV:128:", "avx512vl,avx512bw")
 TARGET_BUILTIN(__builtin_ia32_cmpd128_mask, "UcV4iV4iIiUc", "ncV:128:", "avx512vl")
@@ -1349,7 +1349,7 @@ TARGET_BUILTIN(__builtin_ia32_vpmadd52luq128, "V2OiV2OiV2OiV2Oi", "ncV:128:", "a
 TARGET_BUILTIN(__builtin_ia32_vpmadd52luq256, "V4OiV4OiV4OiV4Oi", "ncV:256:", "avx512ifma,avx512vl|avxifma")
 TARGET_BUILTIN(__builtin_ia32_vcomisd, "iV2dV2dIiIi", "ncV:128:", "avx512f")
 TARGET_BUILTIN(__builtin_ia32_vcomiss, "iV4fV4fIiIi", "ncV:128:", "avx512f")
-TARGET_BUILTIN(__builtin_ia32_kunpckdi, "UOiUOiUOi", "nc", "avx512bw,evex512")
+TARGET_BUILTIN(__builtin_ia32_kunpckdi, "UOiUOiUOi", "nc", "avx512bw")
 TARGET_BUILTIN(__builtin_ia32_kunpcksi, "UiUiUi", "nc", "avx512bw")
 TARGET_BUILTIN(__builtin_ia32_loaddquhi512_mask, "V32sV32sC*V32sUi", "nV:512:", "avx512bw,evex512")
 TARGET_BUILTIN(__builtin_ia32_loaddquqi512_mask, "V64cV64cC*V64cUOi", "nV:512:", "avx512bw,evex512")
@@ -1665,56 +1665,56 @@ TARGET_BUILTIN(__builtin_ia32_fpclassss_mask, "UcV4fIiUc", "ncV:128:", "avx512dq
 TARGET_BUILTIN(__builtin_ia32_kaddqi, "UcUcUc", "nc", "avx512dq")
 TARGET_BUILTIN(__builtin_ia32_kaddhi, "UsUsUs", "nc", "avx512dq")
 TARGET_BUILTIN(__builtin_ia32_kaddsi, "UiUiUi", "nc", "avx512bw")
-TARGET_BUILTIN(__builtin_ia32_kadddi, "UOiUOiUOi", "nc", "avx512bw,evex512")
+TARGET_BUILTIN(__builtin_ia32_kadddi, "UOiUOiUOi", "nc", "avx512bw")
 TARGET_BUILTIN(__builtin_ia32_kandqi, "UcUcUc", "nc", "avx512dq")
 TARGET_BUILTIN(__builtin_ia32_kandhi, "UsUsUs", "nc", "avx512f")
 TARGET_BUILTIN(__builtin_ia32_kandsi, "UiUiUi", "nc", "avx512bw")
-TARGET_BUILTIN(__builtin_ia32_kanddi, "UOiUOiUOi", "nc", "avx512bw,evex512")
+TARGET_BUILTIN(__builtin_ia32_kanddi, "UOiUOiUOi", "nc", "avx512bw")
 TARGET_BUILTIN(__builtin_ia32_kandnqi, "UcUcUc", "nc", "avx512dq")
 TARGET_BUILTIN(__builtin_ia32_kandnhi, "UsUsUs", "nc", "avx512f")
 TARGET_BUILTIN(__builtin_ia32_kandnsi, "UiUiUi", "nc", "avx512bw")
-TARGET_BUILTIN(__builtin_ia32_kandndi, "UOiUOiUOi", "nc", "avx512bw,evex512")
+TARGET_BUILTIN(__builtin_ia32_kandndi, "UOiUOiUOi", "nc", "avx512bw")
 TARGET_BUILTIN(__builtin_ia32_korqi, "UcUcUc", "nc", "avx512dq")
 TARGET_BUILTIN(__builtin_ia32_korhi, "UsUsUs", "nc", "avx512f")
 TARGET_BUILTIN(__builtin_ia32_korsi, "UiUiUi", "nc", "avx512bw")
-TARGET_BUILTIN(__builtin_ia32_kordi, "UOiUOiUOi", "nc", "avx512bw,evex512")
+TARGET_BUILTIN(__builtin_ia32_kordi, "UOiUOiUOi", "nc", "avx512bw")
 TARGET_BUILTIN(__builtin_ia32_kortestcqi, "iUcUc", "nc", "avx512dq")
 TARGET_BUILTIN(__builtin_ia32_kortestzqi, "iUcUc", "nc", "avx512dq")
 TARGET_BUILTIN(__builtin_ia32_kortestchi, "iUsUs", "nc", "avx512f")
 TARGET_BUILTIN(__builtin_ia32_kortestzhi, "iUsUs", "nc", "avx512f")
 TARGET_BUILTIN(__builtin_ia32_kortestcsi, "iUiUi", "nc", "avx512bw")
 TARGET_BUILTIN(__builtin_ia32_kortestzsi, "iUiUi", "nc", "avx512bw")
-TARGET_BUILTIN(__builtin_ia32_kortestcdi, "iUOiUOi", "nc", "avx512bw,evex512")
-TARGET_BUILTIN(__builtin_ia32_kortestzdi, "iUOiUOi", "nc", "avx512bw,evex512")
+TARGET_BUILTIN(__builtin_ia32_kortestcdi, "iUOiUOi", "nc", "avx512bw")
+TARGET_BUILTIN(__builtin_ia32_kortestzdi, "iUOiUOi", "nc", "avx512bw")
 TARGET_BUILTIN(__builtin_ia32_ktestcqi, "iUcUc", "nc", "avx512dq")
 TARGET_BUILTIN(__builtin_ia32_ktestzqi, "iUcUc", "nc", "avx512dq")
 TARGET_BUILTIN(__builtin_ia32_ktestchi, "iUsUs", "nc", "avx512dq")
 TARGET_BUILTIN(__builtin_ia32_ktestzhi, "iUsUs", "nc", "avx512dq")
 TARGET_BUILTIN(__builtin_ia32_ktestcsi, "iUiUi", "nc", "avx512bw")
 TARGET_BUILTIN(__builtin_ia32_ktestzsi, "iUiUi", "nc", "avx512bw")
-TARGET_BUILTIN(__builtin_ia32_ktestcdi, "iUOiUOi", "nc", "avx512bw,evex512")
-TARGET_BUILTIN(__builtin_ia32_ktestzdi, "iUOiUOi", "nc", "avx512bw,evex512")
+TARGET_BUILTIN(__builtin_ia32_ktestcdi, "iUOiUOi", "nc", "avx512bw")
+TARGET_BUILTIN(__builtin_ia32_ktestzdi, "iUOiUOi", "nc", "avx512bw")
 TARGET_BUILTIN(__builtin_ia32_kunpckhi, "UsUsUs", "nc", "avx512f")
 TARGET_BUILTIN(__builtin_ia32_kxnorqi, "UcUcUc", "nc", "avx512dq")
 TARGET_BUILTIN(__builtin_ia32_kxnorhi, "UsUsUs", "nc", "avx512f")
 TARGET_BUILTIN(__builtin_ia32_kxnorsi, "UiUiUi", "nc", "avx512bw")
-TARGET_BUILTIN(__builtin_ia32_kxnordi, "UOiUOiUOi", "nc", "avx512bw,evex512")
+TARGET_BUILTIN(__builtin_ia32_kxnordi, "UOiUOiUOi", "nc", "avx512bw")
 TARGET_BUILTIN(__builtin_ia32_kxorqi, "UcUcUc", "nc", "avx512dq")
 TARGET_BUILTIN(__builtin_ia32_kxorhi, "UsUsUs", "nc", "avx512f")
 TARGET_BUILTIN(__builtin_ia32_kxorsi, "UiUiUi", "nc", "avx512bw")
-TARGET_BUILTIN(__builtin_ia32_kxordi, "UOiUOiUOi", "nc", "avx512bw,evex512")
+TARGET_BUILTIN(__builtin_ia32_kxordi, "UOiUOiUOi", "nc", "avx512bw")
 TARGET_BUILTIN(__builtin_ia32_kshiftliqi, "UcUcIUi", "nc", "avx512dq")
 TARGET_BUILTIN(__builtin_ia32_kshiftlihi, "UsUsIUi", "nc", "avx512f")
 TARGET_BUILTIN(__builtin_ia32_kshiftlisi, "UiUiIUi", "nc", "avx512bw")
-TARGET_BUILTIN(__builtin_ia32_kshiftlidi, "UOiUOiIUi", "nc", "avx512bw,evex512")
+TARGET_BUILTIN(__builtin_ia32_kshiftlidi, "UOiUOiIUi", "nc", "avx512bw")
 TARGET_BUILTIN(__builtin_ia32_kshiftriqi, "UcUcIUi", "nc", "avx512dq")
 TARGET_BUILTIN(__builtin_ia32_kshiftrihi, "UsUsIUi", "nc", "avx512f")
 TARGET_BUILTIN(__builtin_ia32_kshiftrisi, "UiUiIUi", "nc", "avx512bw")
-TARGET_BUILTIN(__builtin_ia32_kshiftridi, "UOiUOiIUi", "nc", "avx512bw,evex512")
+TARGET_BUILTIN(__builtin_ia32_kshiftridi, "UOiUOiIUi", "nc", "avx512bw")
 TARGET_BUILTIN(__builtin_ia32_kmovb, "UcUc", "nc", "avx512dq")
 TARGET_BUILTIN(__builtin_ia32_kmovw, "UsUs", "nc", "avx512f")
 TARGET_BUILTIN(__builtin_ia32_kmovd, "UiUi", "nc", "avx512bw")
-TARGET_BUILTIN(__builtin_ia32_kmovq, "UOiUOi", "nc", "avx512bw,evex512")
+TARGET_BUILTIN(__builtin_ia32_kmovq, "UOiUOi", "nc", "avx512bw")
 TARGET_BUILTIN(__builtin_ia32_palignr512, "V64cV64cV64cIi", "ncV:512:", "avx512bw,evex512")
 TARGET_BUILTIN(__builtin_ia32_dbpsadbw128, "V8sV16cV16cIi", "ncV:128:", "avx512bw,avx512vl")
 TARGET_BUILTIN(__builtin_ia32_dbpsadbw256, "V16sV32cV32cIi", "ncV:256:", "avx512bw,avx512vl")

clang/include/clang/Basic/Cuda.h

@@ -39,9 +39,11 @@ enum class CudaVersion {
   CUDA_118,
   CUDA_120,
   CUDA_121,
-  FULLY_SUPPORTED = CUDA_118,
+  CUDA_122,
+  CUDA_123,
+  FULLY_SUPPORTED = CUDA_123,
   PARTIALLY_SUPPORTED =
-      CUDA_121, // Partially supported. Proceed with a warning.
+      CUDA_123, // Partially supported. Proceed with a warning.
   NEW = 10000,  // Too new. Issue a warning, but allow using it.
 };
 const char *CudaVersionToString(CudaVersion V);
@@ -71,6 +73,7 @@ enum class CudaArch {
   SM_87,
   SM_89,
   SM_90,
+  SM_90a,
   GFX600,
   GFX601,
   GFX602,

clang/include/clang/Basic/DiagnosticGroups.td

@@ -40,6 +40,7 @@ def : DiagGroup<"auto-import">;
 def FrameworkHdrQuotedInclude : DiagGroup<"quoted-include-in-framework-header">;
 def FrameworkIncludePrivateFromPublic :
   DiagGroup<"framework-include-private-from-public">;
+def DeprecatedModuleDotMap : DiagGroup<"deprecated-module-dot-map">;
 def FrameworkHdrAtImport : DiagGroup<"atimport-in-framework-header">;
 def CXX14BinaryLiteral : DiagGroup<"c++14-binary-literal">;
 def CXXPre14CompatBinaryLiteral : DiagGroup<"c++98-c++11-compat-binary-literal">;

clang/include/clang/Basic/DiagnosticLexKinds.td

@@ -884,6 +884,9 @@ def warn_quoted_include_in_framework_header : Warning<
 def warn_framework_include_private_from_public : Warning<
   "public framework header includes private framework header '%0'"
   >, InGroup<FrameworkIncludePrivateFromPublic>;
+def warn_deprecated_module_dot_map : Warning<
+  "'%0' as a module map name is deprecated, rename it to %select{module.modulemap|module.private.modulemap}1%select{| in the 'Modules' directory of the framework}2">,
+  InGroup<DeprecatedModuleDotMap>;
 
 def remark_pp_include_directive_modular_translation : Remark<
   "treating #%select{include|import|include_next|__include_macros}0 as an "

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -1669,6 +1669,8 @@ def err_qualified_friend_def : Error<
   "friend function definition cannot be qualified with '%0'">;
 def err_friend_def_in_local_class : Error<
   "friend function cannot be defined in a local class">;
+def err_friend_specialization_def : Error<
+  "friend function specialization cannot be defined">;
 def err_friend_not_first_in_declaration : Error<
   "'friend' must appear first in a non-function declaration">;
 def err_using_decl_friend : Error<
@@ -8727,7 +8729,7 @@ def err_atomic_op_needs_atomic_int : Error<
   "address argument to atomic operation must be a pointer to "
   "%select{|atomic }0integer (%1 invalid)">;
 def warn_atomic_op_has_invalid_memory_order : Warning<
-  "memory order argument to atomic operation is invalid">,
+  "%select{|success |failure }0memory order argument to atomic operation is invalid">,
   InGroup<DiagGroup<"atomic-memory-ordering">>;
 def err_atomic_op_has_invalid_synch_scope : Error<
   "synchronization scope argument to atomic operation is invalid">;
@@ -12000,7 +12002,7 @@ def warn_tcb_enforcement_violation : Warning<
 
 // RISC-V builtin required extension warning
 def err_riscv_builtin_requires_extension : Error<
-  "builtin requires%select{| at least one of the following extensions to be enabled}0: %1">;
+  "builtin requires%select{| at least one of the following extensions}0: %1">;
 def err_riscv_builtin_invalid_lmul : Error<
   "LMUL argument must be in the range [0,3] or [5,7]">;
 def err_riscv_type_requires_extension : Error<

clang/include/clang/Basic/FPOptions.def

@@ -28,4 +28,5 @@ OPTION(FPEvalMethod, LangOptions::FPEvalMethodKind, 2, AllowApproxFunc)
 OPTION(Float16ExcessPrecision, LangOptions::ExcessPrecisionKind, 2, FPEvalMethod)
 OPTION(BFloat16ExcessPrecision, LangOptions::ExcessPrecisionKind, 2, Float16ExcessPrecision)
 OPTION(MathErrno, bool, 1, BFloat16ExcessPrecision)
+OPTION(ComplexRange, LangOptions::ComplexRangeKind, 2, MathErrno)
 #undef OPTION

clang/include/clang/Basic/Features.def

@@ -104,6 +104,7 @@ FEATURE(scudo, LangOpts.Sanitize.hasOneOf(SanitizerKind::Scudo))
 FEATURE(swiftasynccc,
   PP.getTargetInfo().checkCallingConvention(CC_SwiftAsync) ==
   clang::TargetInfo::CCCR_OK)
+FEATURE(pragma_stdc_cx_limited_range, true)
 // Objective-C features
 FEATURE(objc_arr, LangOpts.ObjCAutoRefCount) // FIXME: REMOVE?
 FEATURE(objc_arc, LangOpts.ObjCAutoRefCount)
@@ -281,7 +282,6 @@ EXTENSION(matrix_types_scalar_division, true)
 EXTENSION(cxx_attributes_on_using_declarations, LangOpts.CPlusPlus11)
 EXTENSION(datasizeof, LangOpts.CPlusPlus)
 
-FEATURE(builtin_headers_in_system_modules, LangOpts.BuiltinHeadersInSystemModules)
 FEATURE(cxx_abi_relative_vtable, LangOpts.CPlusPlus && LangOpts.RelativeCXXABIVTables)
 
 // CUDA/HIP Features

clang/include/clang/Basic/IdentifierTable.h

@@ -511,7 +511,7 @@ class alignas(IdentifierInfoAlignment) IdentifierInfo {
   ///   function(<#int x#>);
   /// \endcode
   bool isEditorPlaceholder() const {
-    return getName().startswith("<#") && getName().endswith("#>");
+    return getName().starts_with("<#") && getName().ends_with("#>");
   }
 
   /// Determine whether \p this is a name reserved for the implementation (C99

clang/include/clang/Basic/LangOptions.def

@@ -220,6 +220,8 @@ BENIGN_LANGOPT(NoSignedZero      , 1, 0, "Permit Floating Point optimization wit
 BENIGN_LANGOPT(AllowRecip        , 1, 0, "Permit Floating Point reciprocal")
 BENIGN_LANGOPT(ApproxFunc        , 1, 0, "Permit Floating Point approximation")
 
+ENUM_LANGOPT(ComplexRange, ComplexRangeKind, 2, CX_Full, "Enable use of range reduction for complex arithmetics.")
+
 BENIGN_LANGOPT(ObjCGCBitmapPrint , 1, 0, "printing of GC's bitmap layout for __weak/__strong ivars")
 
 BENIGN_LANGOPT(AccessControl     , 1, 1, "C++ access control")

clang/include/clang/Basic/LangOptions.h

@@ -392,6 +392,8 @@ class LangOptions : public LangOptionsBase {
     IncompleteOnly = 3,
   };
 
+  enum ComplexRangeKind { CX_Full, CX_Limited, CX_Fortran };
+
 public:
   /// The used language standard.
   LangStandard::Kind LangStd;
@@ -741,6 +743,7 @@ class FPOptions {
       setAllowFEnvAccess(true);
     else
       setAllowFEnvAccess(LangOptions::FPM_Off);
+    setComplexRange(LO.getComplexRange());
   }
 
   bool allowFPContractWithinStatement() const {

clang/include/clang/Basic/TargetBuiltins.h

@@ -309,7 +309,7 @@ namespace clang {
     bool isTupleSet() const { return Flags & IsTupleSet; }
     bool isReadZA() const { return Flags & IsReadZA; }
     bool isWriteZA() const { return Flags & IsWriteZA; }
-
+    bool isReductionQV() const { return Flags & IsReductionQV; }
     uint64_t getBits() const { return Flags; }
     bool isFlagSet(uint64_t Flag) const { return Flags & Flag; }
   };

clang/include/clang/Basic/TargetInfo.h

@@ -266,7 +266,6 @@ class TargetInfo : public TransferrableTargetInfo,
   LLVM_PREFERRED_TYPE(bool)
   unsigned AllowAMDGPUUnsafeFPAtomics : 1;
 
-  LLVM_PREFERRED_TYPE(bool)
   unsigned ARMCDECoprocMask : 8;
 
   unsigned MaxOpenCLWorkGroupSize;
@@ -1424,6 +1423,8 @@ class TargetInfo : public TransferrableTargetInfo,
 
   /// Identify whether this target supports IFuncs.
   bool supportsIFunc() const {
+    if (getTriple().isOSBinFormatMachO())
+      return true;
     return getTriple().isOSBinFormatELF() &&
            ((getTriple().isOSLinux() && !getTriple().isMusl()) ||
             getTriple().isOSFreeBSD());

clang/include/clang/Basic/TokenKinds.def

@@ -911,6 +911,11 @@ PRAGMA_ANNOTATION(pragma_fenv_access_ms)
 // handles them.
 PRAGMA_ANNOTATION(pragma_fenv_round)
 
+// Annotation for #pragma STDC CX_LIMITED_RANGE
+// The lexer produces these so that they only take effect when the parser
+// handles them.
+PRAGMA_ANNOTATION(pragma_cx_limited_range)
+
 // Annotation for #pragma float_control
 // The lexer produces these so that they only take effect when the parser
 // handles them.

clang/include/clang/Basic/arm_sve.td

@@ -1935,17 +1935,41 @@ def SVBGRP   : SInst<"svbgrp[_{d}]",   "ddd", "UcUsUiUl", MergeNone, "aarch64_sv
 def SVBGRP_N : SInst<"svbgrp[_n_{d}]", "dda", "UcUsUiUl", MergeNone, "aarch64_sve_bgrp_x">;
 }
 
+let TargetGuard = "sve2p1|sme" in {
+def SVPSEL_B : SInst<"svpsel_lane_b8",  "PPPm", "Pc", MergeNone, "", [IsStreamingCompatible], []>;
+def SVPSEL_H : SInst<"svpsel_lane_b16", "PPPm", "Ps", MergeNone, "", [IsStreamingCompatible], []>;
+def SVPSEL_S : SInst<"svpsel_lane_b32", "PPPm", "Pi", MergeNone, "", [IsStreamingCompatible], []>;
+def SVPSEL_D : SInst<"svpsel_lane_b64", "PPPm", "Pl", MergeNone, "", [IsStreamingCompatible], []>;
+def SVPSEL_COUNT_ALIAS_B : SInst<"svpsel_lane_c8",  "}}Pm", "Pc", MergeNone, "", [IsStreamingCompatible], []>;
+def SVPSEL_COUNT_ALIAS_H : SInst<"svpsel_lane_c16", "}}Pm", "Ps", MergeNone, "", [IsStreamingCompatible], []>;
+def SVPSEL_COUNT_ALIAS_S : SInst<"svpsel_lane_c32", "}}Pm", "Pi", MergeNone, "", [IsStreamingCompatible], []>;
+def SVPSEL_COUNT_ALIAS_D : SInst<"svpsel_lane_c64", "}}Pm", "Pl", MergeNone, "", [IsStreamingCompatible], []>;
+}
+
+// Standalone sve2.1 builtins
 let TargetGuard = "sve2p1" in {
-def SVFCLAMP   : SInst<"svclamp[_{d}]", "dddd", "hfd", MergeNone, "aarch64_sve_fclamp", [], []>;
+def SVORQV   : SInst<"svorqv[_{d}]", "{Pd", "csilUcUsUiUl", MergeNone, "aarch64_sve_orqv", [IsReductionQV]>;
+def SVEORQV  : SInst<"sveorqv[_{d}]", "{Pd", "csilUcUsUiUl", MergeNone, "aarch64_sve_eorqv", [IsReductionQV]>;
+def SVADDQV  : SInst<"svaddqv[_{d}]", "{Pd", "hfdcsilUcUsUiUl", MergeNone, "aarch64_sve_addqv", [IsReductionQV]>;
+def SVANDQV  : SInst<"svandqv[_{d}]", "{Pd", "csilUcUsUiUl", MergeNone, "aarch64_sve_andqv", [IsReductionQV]>;
+def SVSMAXQV : SInst<"svmaxqv[_{d}]", "{Pd", "csil", MergeNone, "aarch64_sve_smaxqv", [IsReductionQV]>;
+def SVUMAXQV : SInst<"svmaxqv[_{d}]", "{Pd", "UcUsUiUl", MergeNone, "aarch64_sve_umaxqv", [IsReductionQV]>;
+def SVSMINQV : SInst<"svminqv[_{d}]", "{Pd", "csil", MergeNone, "aarch64_sve_sminqv", [IsReductionQV]>;
+def SVUMINQV : SInst<"svminqv[_{d}]", "{Pd", "UcUsUiUl", MergeNone, "aarch64_sve_uminqv", [IsReductionQV]>;
 
-def SVPEXT_SINGLE : SInst<"svpext_lane_{d}", "P}i", "QcQsQiQl", MergeNone, "aarch64_sve_pext", [], [ImmCheck<1, ImmCheck0_3>]>;
-def SVPEXT_X2     : SInst<"svpext_lane_{d}_x2", "2.P}i", "QcQsQiQl", MergeNone, "aarch64_sve_pext_x2", [], [ImmCheck<1, ImmCheck0_1>]>;
+def SVFMAXNMQV: SInst<"svmaxnmqv[_{d}]", "{Pd", "hfd", MergeNone, "aarch64_sve_fmaxnmqv", [IsReductionQV]>;
+def SVFMINNMQV: SInst<"svminnmqv[_{d}]", "{Pd", "hfd", MergeNone, "aarch64_sve_fminnmqv", [IsReductionQV]>;
+def SVFMAXQV: SInst<"svmaxqv[_{d}]", "{Pd", "hfd", MergeNone, "aarch64_sve_fmaxqv", [IsReductionQV]>;
+def SVFMINQV: SInst<"svminqv[_{d}]", "{Pd", "hfd", MergeNone, "aarch64_sve_fminqv", [IsReductionQV]>;
+}
 
-def SVPSEL_COUNT_ALIAS_B : SInst<"svpsel_lane_c8",  "}}Pm", "Pc", MergeNone, "", [], []>;
-def SVPSEL_COUNT_ALIAS_H : SInst<"svpsel_lane_c16", "}}Pm", "Ps", MergeNone, "", [], []>;
-def SVPSEL_COUNT_ALIAS_S : SInst<"svpsel_lane_c32", "}}Pm", "Pi", MergeNone, "", [], []>;
-def SVPSEL_COUNT_ALIAS_D : SInst<"svpsel_lane_c64", "}}Pm", "Pl", MergeNone, "", [], []>;
+let TargetGuard = "sve2p1|sme2" in {
+//FIXME: Replace IsStreamingCompatible with IsStreamingOrHasSVE2p1 when available
+def SVPEXT_SINGLE : SInst<"svpext_lane_{d}", "P}i", "QcQsQiQl", MergeNone, "aarch64_sve_pext", [IsStreamingCompatible], [ImmCheck<1, ImmCheck0_3>]>;
+def SVPEXT_X2     : SInst<"svpext_lane_{d}_x2", "2.P}i", "QcQsQiQl", MergeNone, "aarch64_sve_pext_x2", [IsStreamingCompatible], [ImmCheck<1, ImmCheck0_1>]>;
+}
 
+let TargetGuard = "sve2p1" in {
 def SVWHILEGE_COUNT  : SInst<"svwhilege_{d}",  "}lli", "QcQsQiQl", MergeNone, "aarch64_sve_whilege_{d}", [IsOverloadNone], [ImmCheck<2, ImmCheck2_4_Mul2>]>;
 def SVWHILEGT_COUNT  : SInst<"svwhilegt_{d}",  "}lli", "QcQsQiQl", MergeNone, "aarch64_sve_whilegt_{d}", [IsOverloadNone], [ImmCheck<2, ImmCheck2_4_Mul2>]>;
 def SVWHILELE_COUNT  : SInst<"svwhilele_{d}",  "}lli", "QcQsQiQl", MergeNone, "aarch64_sve_whilele_{d}", [IsOverloadNone], [ImmCheck<2, ImmCheck2_4_Mul2>]>;
@@ -2045,13 +2069,6 @@ let TargetGuard = "sve2p1" in {
 def SVSCLAMP : SInst<"svclamp[_{d}]", "dddd", "csil",     MergeNone, "aarch64_sve_sclamp", [], []>;
 def SVUCLAMP : SInst<"svclamp[_{d}]", "dddd", "UcUsUiUl", MergeNone, "aarch64_sve_uclamp", [], []>;
 
-def SVPSEL_B : SInst<"svpsel_lane_b8",  "PPPm", "Pc", MergeNone, "", [], []>;
-def SVPSEL_H : SInst<"svpsel_lane_b16", "PPPm", "Ps", MergeNone, "", [], []>;
-def SVPSEL_S : SInst<"svpsel_lane_b32", "PPPm", "Pi", MergeNone, "", [], []>;
-def SVPSEL_D : SInst<"svpsel_lane_b64", "PPPm", "Pl", MergeNone, "", [], []>;
-
-def SVCNTP_COUNT : SInst<"svcntp_{d}", "n}i", "QcQsQiQl", MergeNone, "aarch64_sve_cntp_{d}", [IsOverloadNone], [ImmCheck<1, ImmCheck2_4_Mul2>]>;
-
 defm SVREVD : SInstZPZ<"svrevd", "csilUcUsUiUl", "aarch64_sve_revd">;
 }
 
@@ -2060,6 +2077,9 @@ let TargetGuard = "sve2p1|sme2" in {
   def SVPTRUE_COUNT  : SInst<"svptrue_{d}", "}v", "QcQsQiQl", MergeNone, "aarch64_sve_ptrue_{d}", [IsOverloadNone, IsStreamingCompatible], []>;
 
   def SVPFALSE_COUNT_ALIAS : SInst<"svpfalse_c", "}v", "", MergeNone, "", [IsOverloadNone, IsStreamingCompatible]>;
+
+  def SVFCLAMP   : SInst<"svclamp[_{d}]", "dddd", "hfd", MergeNone, "aarch64_sve_fclamp", [IsStreamingCompatible], []>;
+  def SVCNTP_COUNT : SInst<"svcntp_{d}", "n}i", "QcQsQiQl", MergeNone, "aarch64_sve_cntp_{d}", [IsOverloadNone, IsStreamingCompatible], [ImmCheck<1, ImmCheck2_4_Mul2>]>;
 }
 
 let TargetGuard = "sve2p1,b16b16" in {
@@ -2164,6 +2184,21 @@ let TargetGuard = "sme2" in {
 
   def REINTERPRET_SVBOOL_TO_SVCOUNT : Inst<"svreinterpret[_c]", "}P", "Pc", MergeNone, "", [IsStreamingCompatible], []>;
   def REINTERPRET_SVCOUNT_TO_SVBOOL : Inst<"svreinterpret[_b]", "P}", "Pc", MergeNone, "", [IsStreamingCompatible], []>;
+
+  // SQDMULH
+  def SVSQDMULH_SINGLE_X2 : SInst<"svqdmulh[_single_{d}_x2]", "22d", "csil", MergeNone, "aarch64_sve_sqdmulh_single_vgx2", [IsStreaming], []>;
+  def SVSQDMULH_SINGLE_X4 : SInst<"svqdmulh[_single_{d}_x4]", "44d", "csil", MergeNone, "aarch64_sve_sqdmulh_single_vgx4", [IsStreaming], []>;
+  def SVSQDMULH_X2        : SInst<"svqdmulh[_{d}_x2]",        "222", "csil", MergeNone, "aarch64_sve_sqdmulh_vgx2",        [IsStreaming], []>;
+  def SVSQDMULH_X4        : SInst<"svqdmulh[_{d}_x4]",        "444", "csil", MergeNone, "aarch64_sve_sqdmulh_vgx4",        [IsStreaming], []>;
+}
+
+let TargetGuard = "sve2p1|sme2" in {
+  // SQRSHRN / UQRSHRN
+  def SVQRSHRN_X2   : SInst<"svqrshrn[_n]_{0}[_{d}_x2]", "h2i", "i",    MergeNone, "aarch64_sve_sqrshrn_x2", [IsStreamingCompatible], [ImmCheck<1, ImmCheck1_16>]>;
+  def SVUQRSHRN_X2  : SInst<"svqrshrn[_n]_{0}[_{d}_x2]", "e2i", "Ui",   MergeNone, "aarch64_sve_uqrshrn_x2", [IsStreamingCompatible], [ImmCheck<1, ImmCheck1_16>]>;
+
+  // SQRSHRUN
+  def SVSQRSHRUN_X2 : SInst<"svqrshrun[_n]_{0}[_{d}_x2]", "e2i", "i",  MergeNone, "aarch64_sve_sqrshrun_x2", [IsStreamingCompatible], [ImmCheck<1, ImmCheck1_16>]>;
 }
 
 let TargetGuard = "sve2p1" in {
@@ -2245,11 +2280,13 @@ let TargetGuard = "sme2" in {
 //
 // Multi-vector saturating extract narrow and interleave
 //
-let TargetGuard = "sme2" in {
+let TargetGuard = "sme2|sve2p1" in {
   def SVQCVTN_S16_S32_X2 : SInst<"svqcvtn_s16[_{d}_x2]", "h2.d", "i", MergeNone, "aarch64_sve_sqcvtn_x2", [IsStreamingCompatible], []>;
   def SVQCVTN_U16_U32_X2 : SInst<"svqcvtn_u16[_{d}_x2]", "e2.d", "Ui", MergeNone, "aarch64_sve_uqcvtn_x2", [IsStreamingCompatible], []>;
   def SVQCVTN_U16_S32_X2 : SInst<"svqcvtn_u16[_{d}_x2]", "e2.d", "i", MergeNone, "aarch64_sve_sqcvtun_x2", [IsStreamingCompatible], []>;
+}
 
+let TargetGuard = "sme2" in {
   def SVQCVTN_S8_S32_X4 : SInst<"svqcvtn_s8[_{d}_x4]", "q4.d", "i", MergeNone, "aarch64_sve_sqcvtn_x4", [IsStreaming], []>;
   def SVQCVTN_U8_U32_X4 : SInst<"svqcvtn_u8[_{d}_x4]", "b4.d", "Ui", MergeNone, "aarch64_sve_uqcvtn_x4", [IsStreaming], []>;
   def SVQCVTN_U8_S32_X4 : SInst<"svqcvtn_u8[_{d}_x4]", "b4.d", "i", MergeNone, "aarch64_sve_sqcvtun_x4", [IsStreaming], []>;
@@ -2258,3 +2295,14 @@ let TargetGuard = "sme2" in {
   def SVQCVTN_U16_U64_X4 : SInst<"svqcvtn_u16[_{d}_x4]", "b4.d", "Ul", MergeNone, "aarch64_sve_uqcvtn_x4", [IsStreaming], []>;
   def SVQCVTN_U16_S64_X4 : SInst<"svqcvtn_u16[_{d}_x4]", "b4.d", "l", MergeNone, "aarch64_sve_sqcvtun_x4", [IsStreaming], []>;
 }
+
+//
+// Multi-vector unpack
+//
+
+let TargetGuard = "sme2" in {
+  def SVSUNPK_X2 : SInst<"svunpk_{d}[_{1}_x2]", "2h",   "sil",    MergeNone, "aarch64_sve_sunpk_x2", [IsStreaming], []>;
+  def SVUUNPK_X2 : SInst<"svunpk_{d}[_{1}_x2]", "2h",   "UsUiUl", MergeNone, "aarch64_sve_uunpk_x2", [IsStreaming], []>;
+  def SVSUNPK_X4 : SInst<"svunpk_{d}[_{3}_x4]", "42.h", "sil",    MergeNone, "aarch64_sve_sunpk_x4", [IsStreaming], []>;
+  def SVUUNPK_X4 : SInst<"svunpk_{d}[_{3}_x4]", "42.h", "UsUiUl", MergeNone, "aarch64_sve_uunpk_x4", [IsStreaming], []>;
+}

clang/include/clang/Basic/arm_sve_sme_incl.td

@@ -129,6 +129,7 @@
 // Z: const pointer to uint64_t
 
 // Prototype modifiers added for SVE2p1
+// {: 128b vector
 // }: svcount_t
 
 class MergeType<int val, string suffix=""> {
@@ -225,6 +226,7 @@ def IsSharedZA                : FlagType<0x8000000000>;
 def IsPreservesZA             : FlagType<0x10000000000>;
 def IsReadZA                  : FlagType<0x20000000000>;
 def IsWriteZA                 : FlagType<0x40000000000>;
+def IsReductionQV             : FlagType<0x80000000000>;
 
 // These must be kept in sync with the flags in include/clang/Basic/TargetBuiltins.h
 class ImmCheckType<int val> {

clang/include/clang/Driver/Options.td

@@ -1010,6 +1010,30 @@ defm offload_uniform_block : BoolFOption<"offload-uniform-block",
   NegFlag<SetFalse, [], [ClangOption, CC1Option], "Don't assume">,
   BothFlags<[], [ClangOption], " that kernels are launched with uniform block sizes (default true for CUDA/HIP and false otherwise)">>;
 
+def fcx_limited_range : Joined<["-"], "fcx-limited-range">,
+  Group<f_Group>, Visibility<[ClangOption, CC1Option]>,
+  HelpText<"Basic algebraic expansions of complex arithmetic operations "
+           "involving are enabled.">;
+
+def fno_cx_limited_range : Joined<["-"], "fno-cx-limited-range">,
+  Group<f_Group>, Visibility<[ClangOption, CC1Option]>,
+  HelpText<"Basic algebraic expansions of complex arithmetic operations "
+           "involving are disabled.">;
+
+def fcx_fortran_rules : Joined<["-"], "fcx-fortran-rules">,
+  Group<f_Group>, Visibility<[ClangOption, CC1Option]>,
+  HelpText<"Range reduction is enabled for complex arithmetic operations.">;
+
+def fno_cx_fortran_rules : Joined<["-"], "fno-cx-fortran-rules">,
+  Group<f_Group>, Visibility<[ClangOption, CC1Option]>,
+  HelpText<"Range reduction is disabled for complex arithmetic operations.">;
+
+def complex_range_EQ : Joined<["-"], "complex-range=">, Group<f_Group>,
+  Visibility<[CC1Option]>,
+  Values<"full,limited,fortran">, NormalizedValuesScope<"LangOptions">,
+  NormalizedValues<["CX_Full", "CX_Limited", "CX_Fortran"]>,
+  MarshallingInfoEnum<LangOpts<"ComplexRange">, "CX_Full">;
+
 // OpenCL-only Options
 def cl_opt_disable : Flag<["-"], "cl-opt-disable">, Group<opencl_Group>,
   Visibility<[ClangOption, CC1Option]>,
@@ -2872,9 +2896,11 @@ defm asm_blocks : BoolFOption<"asm-blocks",
   LangOpts<"AsmBlocks">, Default<fms_extensions.KeyPath>,
   PosFlag<SetTrue, [], [ClangOption, CC1Option]>,
   NegFlag<SetFalse>>;
-def fms_volatile : Flag<["-"], "fms-volatile">, Group<f_Group>,
-  Visibility<[ClangOption, CC1Option]>,
-  MarshallingInfoFlag<LangOpts<"MSVolatile">>;
+defm ms_volatile : BoolFOption<"ms-volatile",
+  LangOpts<"MSVolatile">, DefaultFalse,
+  PosFlag<SetTrue, [], [ClangOption, CC1Option],
+          "Volatile loads and stores have acquire and release semantics">,
+  NegFlag<SetFalse>>;
 def fmsc_version : Joined<["-"], "fmsc-version=">, Group<f_Group>,
   Visibility<[ClangOption, CLOption]>,
   HelpText<"Microsoft compiler version number to report in _MSC_VER (0 = don't define it (default))">;
@@ -6354,6 +6380,12 @@ def J : JoinedOrSeparate<["-"], "J">,
   Group<gfortran_Group>,
   Alias<module_dir>;
 
+let Visibility = [FlangOption] in {
+def no_fortran_main : Flag<["-"], "fno-fortran-main">,
+  Visibility<[FlangOption]>, Group<f_Group>,
+  HelpText<"Do not include Fortran_main.a (provided by Flang) when linking">;
+} // let Visibility = [ FlangOption ]
+
 //===----------------------------------------------------------------------===//
 // FC1 Options
 //===----------------------------------------------------------------------===//
@@ -8187,7 +8219,7 @@ def _SLASH_winsysroot : CLJoinedOrSeparate<"winsysroot">,
   HelpText<"Same as \"/diasdkdir <dir>/DIA SDK\" /vctoolsdir <dir>/VC/Tools/MSVC/<vctoolsversion> \"/winsdkdir <dir>/Windows Kits/10\"">,
   MetaVarName<"<dir>">;
 def _SLASH_volatile_iso : Option<["/", "-"], "volatile:iso", KIND_FLAG>,
-  Group<_SLASH_volatile_Group>, Flags<[NoXarchOption]>, Visibility<[CLOption]>,
+  Visibility<[CLOption]>, Alias<fno_ms_volatile>,
   HelpText<"Volatile loads and stores have standard semantics">;
 def _SLASH_vmb : CLFlag<"vmb">,
   HelpText<"Use a best-case representation method for member pointers">;
@@ -8202,7 +8234,7 @@ def _SLASH_vmv : CLFlag<"vmv">,
   HelpText<"Set the default most-general representation to "
            "virtual inheritance">;
 def _SLASH_volatile_ms  : Option<["/", "-"], "volatile:ms", KIND_FLAG>,
-  Group<_SLASH_volatile_Group>, Flags<[NoXarchOption]>, Visibility<[CLOption]>,
+  Visibility<[CLOption]>, Alias<fms_volatile>,
   HelpText<"Volatile loads and stores have acquire and release semantics">;
 def _SLASH_clang : CLJoined<"clang:">,
   HelpText<"Pass <arg> to the clang driver">, MetaVarName<"<arg>">;

clang/include/clang/Parse/Parser.h

@@ -769,6 +769,10 @@ class Parser : public CodeCompletionHandler {
   /// #pragma STDC FENV_ROUND...
   void HandlePragmaFEnvRound();
 
+  /// Handle the annotation token produced for
+  /// #pragma STDC CX_LIMITED_RANGE...
+  void HandlePragmaCXLimitedRange();
+
   /// Handle the annotation token produced for
   /// #pragma float_control
   void HandlePragmaFloatControl();

clang/include/clang/Sema/Sema.h

@@ -2102,7 +2102,7 @@ class Sema final {
                                  SourceLocation AttrLoc);
 
   CodeAlignAttr *BuildCodeAlignAttr(const AttributeCommonInfo &CI, Expr *E);
-  bool CheckRebuiltCodeAlignStmtAttributes(ArrayRef<const Attr *> Attrs);
+  bool CheckRebuiltStmtAttributes(ArrayRef<const Attr *> Attrs);
 
   bool CheckQualifiedFunctionForTypeId(QualType T, SourceLocation Loc);
 
@@ -7311,8 +7311,7 @@ class Sema final {
 
   /// ActOnLambdaExpr - This is called when the body of a lambda expression
   /// was successfully completed.
-  ExprResult ActOnLambdaExpr(SourceLocation StartLoc, Stmt *Body,
-                             Scope *CurScope);
+  ExprResult ActOnLambdaExpr(SourceLocation StartLoc, Stmt *Body);
 
   /// Does copying/destroying the captured variable have side effects?
   bool CaptureHasSideEffects(const sema::Capture &From);
@@ -11023,6 +11022,11 @@ class Sema final {
   /// \#pragma STDC FENV_ACCESS
   void ActOnPragmaFEnvAccess(SourceLocation Loc, bool IsEnabled);
 
+  /// ActOnPragmaCXLimitedRange - Called on well formed
+  /// \#pragma STDC CX_LIMITED_RANGE
+  void ActOnPragmaCXLimitedRange(SourceLocation Loc,
+                                 LangOptions::ComplexRangeKind Range);
+
   /// Called on well formed '\#pragma clang fp' that has option 'exceptions'.
   void ActOnPragmaFPExceptions(SourceLocation Loc,
                                LangOptions::FPExceptionModeKind);

clang/include/clang/Serialization/ASTReader.h

@@ -1424,7 +1424,7 @@ class ASTReader
   RecordLocation TypeCursorForIndex(unsigned Index);
   void LoadedDecl(unsigned Index, Decl *D);
   Decl *ReadDeclRecord(serialization::DeclID ID);
-  void markIncompleteDeclChain(Decl *Canon);
+  void markIncompleteDeclChain(Decl *D);
 
   /// Returns the most recent declaration of a declaration (which must be
   /// of a redeclarable kind) that is either local or has already been loaded
@@ -2093,7 +2093,7 @@ class ASTReader
            SmallVectorImpl<std::pair<Selector, SourceLocation>> &Sels) override;
 
   void ReadWeakUndeclaredIdentifiers(
-           SmallVectorImpl<std::pair<IdentifierInfo *, WeakInfo>> &WI) override;
+      SmallVectorImpl<std::pair<IdentifierInfo *, WeakInfo>> &WeakIDs) override;
 
   void ReadUsedVTables(SmallVectorImpl<ExternalVTableUse> &VTables) override;
 
@@ -2203,7 +2203,7 @@ class ASTReader
 
   /// Retrieve the global selector ID that corresponds to this
   /// the local selector ID in a given module.
-  serialization::SelectorID getGlobalSelectorID(ModuleFile &F,
+  serialization::SelectorID getGlobalSelectorID(ModuleFile &M,
                                                 unsigned LocalID) const;
 
   /// Read the contents of a CXXCtorInitializer array.
@@ -2415,12 +2415,7 @@ class BitsUnpacker {
   BitsUnpacker(BitsUnpacker &&) = delete;
   BitsUnpacker operator=(const BitsUnpacker &) = delete;
   BitsUnpacker operator=(BitsUnpacker &&) = delete;
-  ~BitsUnpacker() {
-#ifndef NDEBUG
-    while (isValid())
-      assert(!getNextBit() && "There are unprocessed bits!");
-#endif
-  }
+  ~BitsUnpacker() = default;
 
   void updateValue(uint32_t V) {
     Value = V;

clang/include/clang/StaticAnalyzer/Checkers/Checkers.td

@@ -36,6 +36,7 @@ def CoreAlpha : Package<"core">, ParentPackage<Alpha>;
 // Note: OptIn is *not* intended for checkers that are too noisy to be on by
 // default. Such checkers belong in the alpha package.
 def OptIn : Package<"optin">;
+def CoreOptIn : Package<"core">, ParentPackage<OptIn>;
 
 // In the Portability package reside checkers for finding code that relies on
 // implementation-defined behavior. Such checks are wanted for cross-platform
@@ -439,6 +440,18 @@ def UndefinedNewArraySizeChecker : Checker<"NewArraySize">,
 
 } // end "core.uninitialized"
 
+//===----------------------------------------------------------------------===//
+// Optin checkers for core language features
+//===----------------------------------------------------------------------===//
+
+let ParentPackage = CoreOptIn in {
+
+def EnumCastOutOfRangeChecker : Checker<"EnumCastOutOfRange">,
+  HelpText<"Check integer to enumeration casts for out of range values">,
+  Documentation<HasDocumentation>;
+
+} // end "optin.core"
+
 //===----------------------------------------------------------------------===//
 // Unix API checkers.
 //===----------------------------------------------------------------------===//
@@ -774,10 +787,6 @@ def DeleteWithNonVirtualDtorChecker : Checker<"DeleteWithNonVirtualDtor">,
            "destructor in their base class">,
   Documentation<HasDocumentation>;
 
-def EnumCastOutOfRangeChecker : Checker<"EnumCastOutOfRange">,
-  HelpText<"Check integer to enumeration casts for out of range values">,
-  Documentation<HasDocumentation>;
-
 def IteratorModeling : Checker<"IteratorModeling">,
   HelpText<"Models iterators of C++ containers">,
   Dependencies<[ContainerModeling]>,

clang/include/clang/StaticAnalyzer/Core/AnalyzerOptions.h

@@ -409,8 +409,8 @@ AnalyzerOptions::getRegisteredCheckers(bool IncludeExperimental) {
   };
   std::vector<StringRef> Checkers;
   for (StringRef CheckerName : StaticAnalyzerCheckerNames) {
-    if (!CheckerName.startswith("debug.") &&
-        (IncludeExperimental || !CheckerName.startswith("alpha.")))
+    if (!CheckerName.starts_with("debug.") &&
+        (IncludeExperimental || !CheckerName.starts_with("alpha.")))
       Checkers.push_back(CheckerName);
   }
   return Checkers;

clang/include/clang/StaticAnalyzer/Core/BugReporter/BugReporter.h

@@ -19,6 +19,7 @@
 #include "clang/Basic/SourceLocation.h"
 #include "clang/Lex/Preprocessor.h"
 #include "clang/StaticAnalyzer/Core/BugReporter/BugReporterVisitors.h"
+#include "clang/StaticAnalyzer/Core/BugReporter/BugSuppression.h"
 #include "clang/StaticAnalyzer/Core/BugReporter/BugType.h"
 #include "clang/StaticAnalyzer/Core/CheckerManager.h"
 #include "clang/StaticAnalyzer/Core/PathSensitive/ExplodedGraph.h"
@@ -594,6 +595,9 @@ class BugReporter {
   /// A vector of BugReports for tracking the allocated pointers and cleanup.
   std::vector<BugReportEquivClass *> EQClassesVector;
 
+  /// User-provided in-code suppressions.
+  BugSuppression UserSuppressions;
+
 public:
   BugReporter(BugReporterData &d);
   virtual ~BugReporter();

clang/include/clang/StaticAnalyzer/Core/BugReporter/BugSuppression.h

@@ -0,0 +1,53 @@
+//===- BugSuppression.h - Suppression interface -----------------*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+//  This file defines BugSuppression, a simple interface class encapsulating
+//  all user provided in-code suppressions.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_CLANG_STATICANALYZER_CORE_BUGREPORTER_SUPPRESSION_H
+#define LLVM_CLANG_STATICANALYZER_CORE_BUGREPORTER_SUPPRESSION_H
+
+#include "clang/Basic/SourceLocation.h"
+#include "llvm/ADT/DenseMap.h"
+#include "llvm/ADT/SmallVector.h"
+
+namespace clang {
+class Decl;
+
+namespace ento {
+class BugReport;
+class PathDiagnosticLocation;
+
+class BugSuppression {
+public:
+  using DiagnosticIdentifierList = llvm::ArrayRef<llvm::StringRef>;
+
+  /// Return true if the given bug report was explicitly suppressed by the user.
+  bool isSuppressed(const BugReport &);
+
+  /// Return true if the bug reported at the given location was explicitly
+  /// suppressed by the user.
+  bool isSuppressed(const PathDiagnosticLocation &Location,
+                    const Decl *DeclWithIssue,
+                    DiagnosticIdentifierList DiagnosticIdentification);
+
+private:
+  // Overly pessimistic number, to be honest.
+  static constexpr unsigned EXPECTED_NUMBER_OF_SUPPRESSIONS = 8;
+  using CachedRanges =
+      llvm::SmallVector<SourceRange, EXPECTED_NUMBER_OF_SUPPRESSIONS>;
+
+  llvm::DenseMap<const Decl *, CachedRanges> CachedSuppressionLocations;
+};
+
+} // end namespace ento
+} // end namespace clang
+
+#endif // LLVM_CLANG_STATICANALYZER_CORE_BUGREPORTER_SUPPRESSION_H

clang/lib/APINotes/APINotesManager.cpp

@@ -198,7 +198,7 @@ static void checkPrivateAPINotesName(DiagnosticsEngine &Diags,
   StringRef RealFileName =
       llvm::sys::path::filename(File->tryGetRealPathName());
   StringRef RealStem = llvm::sys::path::stem(RealFileName);
-  if (RealStem.endswith("_private"))
+  if (RealStem.ends_with("_private"))
     return;
 
   unsigned DiagID = diag::warn_apinotes_private_case;

clang/lib/APINotes/APINotesYAMLCompiler.cpp

@@ -745,7 +745,7 @@ class YAMLConverter {
     convertCommonEntity(M, MI, M.Selector);
 
     // Check if the selector ends with ':' to determine if it takes arguments.
-    bool takesArguments = M.Selector.endswith(":");
+    bool takesArguments = M.Selector.ends_with(":");
 
     // Split the selector into pieces.
     llvm::SmallVector<StringRef, 4> Args;

clang/lib/ARCMigrate/ARCMT.cpp

@@ -201,7 +201,7 @@ createInvocationForMigration(CompilerInvocation &origCI,
   for (std::vector<std::string>::iterator
          I = CInvok->getDiagnosticOpts().Warnings.begin(),
          E = CInvok->getDiagnosticOpts().Warnings.end(); I != E; ++I) {
-    if (!StringRef(*I).startswith("error"))
+    if (!StringRef(*I).starts_with("error"))
       WarnOpts.push_back(*I);
   }
   WarnOpts.push_back("error=arc-unsafe-retained-assign");

clang/lib/ARCMigrate/ObjCMT.cpp

@@ -562,7 +562,7 @@ static void rewriteToObjCProperty(const ObjCMethodDecl *Getter,
 static bool IsCategoryNameWithDeprecatedSuffix(ObjCContainerDecl *D) {
   if (ObjCCategoryDecl *CatDecl = dyn_cast<ObjCCategoryDecl>(D)) {
     StringRef Name = CatDecl->getName();
-    return Name.endswith("Deprecated");
+    return Name.ends_with("Deprecated");
   }
   return false;
 }
@@ -1176,12 +1176,12 @@ bool ObjCMigrateASTConsumer::migrateProperty(ASTContext &Ctx,
   if (!SetterMethod) {
     // try a different naming convention for getter: isXxxxx
     StringRef getterNameString = getterName->getName();
-    bool IsPrefix = getterNameString.startswith("is");
+    bool IsPrefix = getterNameString.starts_with("is");
     // Note that we don't want to change an isXXX method of retainable object
     // type to property (readonly or otherwise).
     if (IsPrefix && GRT->isObjCRetainableType())
       return false;
-    if (IsPrefix || getterNameString.startswith("get")) {
+    if (IsPrefix || getterNameString.starts_with("get")) {
       LengthOfPrefix = (IsPrefix ? 2 : 3);
       const char *CGetterName = getterNameString.data() + LengthOfPrefix;
       // Make sure that first character after "is" or "get" prefix can
@@ -1320,11 +1320,11 @@ void ObjCMigrateASTConsumer::migrateFactoryMethod(ASTContext &Ctx,
   if (OIT_Family == OIT_Singleton || OIT_Family == OIT_ReturnsSelf) {
     StringRef STRefMethodName(MethodName);
     size_t len = 0;
-    if (STRefMethodName.startswith("standard"))
+    if (STRefMethodName.starts_with("standard"))
       len = strlen("standard");
-    else if (STRefMethodName.startswith("shared"))
+    else if (STRefMethodName.starts_with("shared"))
       len = strlen("shared");
-    else if (STRefMethodName.startswith("default"))
+    else if (STRefMethodName.starts_with("default"))
       len = strlen("default");
     else
       return;
@@ -1341,7 +1341,7 @@ void ObjCMigrateASTConsumer::migrateFactoryMethod(ASTContext &Ctx,
   StringRef LoweredMethodName(MethodName);
   std::string StringLoweredMethodName = LoweredMethodName.lower();
   LoweredMethodName = StringLoweredMethodName;
-  if (!LoweredMethodName.startswith(ClassNamePostfix))
+  if (!LoweredMethodName.starts_with(ClassNamePostfix))
     return;
   if (OIT_Family == OIT_ReturnsSelf)
     ReplaceWithClasstype(*this, OM);

clang/lib/ARCMigrate/TransUnbridgedCasts.cpp

@@ -146,7 +146,7 @@ class UnbridgedCastRewriter : public RecursiveASTVisitor<UnbridgedCastRewriter>{
             ento::cocoa::isRefType(E->getSubExpr()->getType(), "CF",
                                    FD->getIdentifier()->getName())) {
           StringRef fname = FD->getIdentifier()->getName();
-          if (fname.endswith("Retain") || fname.contains("Create") ||
+          if (fname.ends_with("Retain") || fname.contains("Create") ||
               fname.contains("Copy")) {
             // Do not migrate to couple of bridge transfer casts which
             // cancel each other out. Leave it unchanged so error gets user

clang/lib/ARCMigrate/TransformActions.cpp

@@ -431,7 +431,7 @@ bool TransformActionsImpl::canReplaceText(SourceLocation loc, StringRef text) {
   if (invalidTemp)
     return false;
 
-  return file.substr(locInfo.second).startswith(text);
+  return file.substr(locInfo.second).starts_with(text);
 }
 
 void TransformActionsImpl::commitInsert(SourceLocation loc, StringRef text) {

clang/lib/ARCMigrate/Transforms.cpp

@@ -95,7 +95,7 @@ bool trans::isPlusOne(const Expr *E) {
           ento::cocoa::isRefType(callE->getType(), "CF",
                                  FD->getIdentifier()->getName())) {
         StringRef fname = FD->getIdentifier()->getName();
-        if (fname.endswith("Retain") || fname.contains("Create") ||
+        if (fname.ends_with("Retain") || fname.contains("Create") ||
             fname.contains("Copy"))
           return true;
       }

clang/lib/AST/ASTContext.cpp

@@ -8223,7 +8223,7 @@ void ASTContext::getObjCEncodingForTypeImpl(QualType T, std::string &S,
       // Another legacy compatibility encoding. Some ObjC qualifier and type
       // combinations need to be rearranged.
       // Rewrite "in const" from "nr" to "rn"
-      if (StringRef(S).endswith("nr"))
+      if (StringRef(S).ends_with("nr"))
         S.replace(S.end()-2, S.end(), "rn");
     }
 
@@ -13519,7 +13519,7 @@ void ASTContext::getFunctionFeatureMap(llvm::StringMap<bool> &FeatureMap,
                       Target->getTargetOpts().FeaturesAsWritten.begin(),
                       Target->getTargetOpts().FeaturesAsWritten.end());
     } else {
-      if (VersionStr.startswith("arch="))
+      if (VersionStr.starts_with("arch="))
         TargetCPU = VersionStr.drop_front(sizeof("arch=") - 1);
       else if (VersionStr != "default")
         Features.push_back((StringRef{"+"} + VersionStr).str());

clang/lib/AST/Decl.cpp

@@ -4150,6 +4150,7 @@ FunctionDecl::setFunctionTemplateSpecialization(ASTContext &C,
   assert(TSK != TSK_Undeclared &&
          "Must specify the type of function template specialization");
   assert((TemplateOrSpecialization.isNull() ||
+          getFriendObjectKind() != FOK_None ||
           TSK == TSK_ExplicitSpecialization) &&
          "Member specialization must be an explicit specialization");
   FunctionTemplateSpecializationInfo *Info =

clang/lib/AST/DeclPrinter.cpp

@@ -1728,7 +1728,7 @@ void DeclPrinter::VisitObjCPropertyDecl(ObjCPropertyDecl *PDecl) {
   std::string TypeStr = PDecl->getASTContext().getUnqualifiedObjCPointerType(T).
       getAsString(Policy);
   Out << ' ' << TypeStr;
-  if (!StringRef(TypeStr).endswith("*"))
+  if (!StringRef(TypeStr).ends_with("*"))
     Out << ' ';
   Out << *PDecl;
   if (Policy.PolishForDeclaration)

clang/lib/AST/ExprConstant.cpp

@@ -8514,14 +8514,24 @@ bool LValueExprEvaluator::VisitVarDecl(const Expr *E, const VarDecl *VD) {
       return false;
 
     if (auto *FD = Info.CurrentCall->LambdaCaptureFields.lookup(VD)) {
+      const auto *MD = cast<CXXMethodDecl>(Info.CurrentCall->Callee);
+
+      // Static lambda function call operators can't have captures. We already
+      // diagnosed this, so bail out here.
+      if (MD->isStatic()) {
+        assert(Info.CurrentCall->This == nullptr &&
+               "This should not be set for a static call operator");
+        return false;
+      }
+
       // Start with 'Result' referring to the complete closure object...
-      if (auto *MD = cast<CXXMethodDecl>(Info.CurrentCall->Callee);
-          MD->isExplicitObjectMemberFunction()) {
+      if (MD->isExplicitObjectMemberFunction()) {
         APValue *RefValue =
             Info.getParamSlot(Info.CurrentCall->Arguments, MD->getParamDecl(0));
         Result.setFrom(Info.Ctx, *RefValue);
       } else
         Result = *Info.CurrentCall->This;
+
       // ... then update it to refer to the field of the closure object
       // that represents the capture.
       if (!HandleLValueMember(Info, E, Result, FD))

clang/lib/AST/Interp/ByteCodeEmitter.cpp

@@ -61,6 +61,11 @@ ByteCodeEmitter::compileFunc(const FunctionDecl *FuncDecl) {
       MD->getParent()->getCaptureFields(LC, LTC);
 
       for (auto Cap : LC) {
+        // Static lambdas cannot have any captures. If this one does,
+        // it has already been diagnosed and we can only ignore it.
+        if (MD->isStatic())
+          return nullptr;
+
         unsigned Offset = R->getField(Cap.second)->Offset;
         this->LambdaCaptures[Cap.first] = {
             Offset, Cap.second->getType()->isReferenceType()};

clang/lib/AST/Interp/ByteCodeExprGen.cpp

@@ -222,6 +222,64 @@ bool ByteCodeExprGen<Emitter>::VisitCastExpr(const CastExpr *CE) {
     return this->emitNE(PtrT, CE);
   }
 
+  case CK_IntegralComplexToBoolean:
+  case CK_FloatingComplexToBoolean: {
+    std::optional<PrimType> ElemT =
+        classifyComplexElementType(SubExpr->getType());
+    if (!ElemT)
+      return false;
+    // We emit the expression (__real(E) != 0 || __imag(E) != 0)
+    // for us, that means (bool)E[0] || (bool)E[1]
+    if (!this->visit(SubExpr))
+      return false;
+    if (!this->emitConstUint8(0, CE))
+      return false;
+    if (!this->emitArrayElemPtrUint8(CE))
+      return false;
+    if (!this->emitLoadPop(*ElemT, CE))
+      return false;
+    if (*ElemT == PT_Float) {
+      if (!this->emitCastFloatingIntegral(PT_Bool, CE))
+        return false;
+    } else {
+      if (!this->emitCast(*ElemT, PT_Bool, CE))
+        return false;
+    }
+
+    // We now have the bool value of E[0] on the stack.
+    LabelTy LabelTrue = this->getLabel();
+    if (!this->jumpTrue(LabelTrue))
+      return false;
+
+    if (!this->emitConstUint8(1, CE))
+      return false;
+    if (!this->emitArrayElemPtrPopUint8(CE))
+      return false;
+    if (!this->emitLoadPop(*ElemT, CE))
+      return false;
+    if (*ElemT == PT_Float) {
+      if (!this->emitCastFloatingIntegral(PT_Bool, CE))
+        return false;
+    } else {
+      if (!this->emitCast(*ElemT, PT_Bool, CE))
+        return false;
+    }
+    // Leave the boolean value of E[1] on the stack.
+    LabelTy EndLabel = this->getLabel();
+    this->jump(EndLabel);
+
+    this->emitLabel(LabelTrue);
+    if (!this->emitPopPtr(CE))
+      return false;
+    if (!this->emitConstBool(true, CE))
+      return false;
+
+    this->fallthrough(EndLabel);
+    this->emitLabel(EndLabel);
+
+    return true;
+  }
+
   case CK_ToVoid:
     return discard(SubExpr);
 
@@ -258,6 +316,9 @@ bool ByteCodeExprGen<Emitter>::VisitBinaryOperator(const BinaryOperator *BO) {
   if (BO->isLogicalOp())
     return this->VisitLogicalBinOp(BO);
 
+  if (BO->getType()->isAnyComplexType())
+    return this->VisitComplexBinOp(BO);
+
   const Expr *LHS = BO->getLHS();
   const Expr *RHS = BO->getRHS();
 
@@ -500,6 +561,107 @@ bool ByteCodeExprGen<Emitter>::VisitLogicalBinOp(const BinaryOperator *E) {
   return true;
 }
 
+template <class Emitter>
+bool ByteCodeExprGen<Emitter>::VisitComplexBinOp(const BinaryOperator *E) {
+  assert(Initializing);
+
+  const Expr *LHS = E->getLHS();
+  const Expr *RHS = E->getRHS();
+  PrimType LHSElemT = *this->classifyComplexElementType(LHS->getType());
+  PrimType RHSElemT = *this->classifyComplexElementType(RHS->getType());
+
+  unsigned LHSOffset = this->allocateLocalPrimitive(LHS, PT_Ptr, true, false);
+  unsigned RHSOffset = this->allocateLocalPrimitive(RHS, PT_Ptr, true, false);
+  unsigned ResultOffset = ~0u;
+  if (!this->DiscardResult)
+    ResultOffset = this->allocateLocalPrimitive(E, PT_Ptr, true, false);
+
+  assert(LHSElemT == RHSElemT);
+
+  // Save result pointer in ResultOffset
+  if (!this->DiscardResult) {
+    if (!this->emitDupPtr(E))
+      return false;
+    if (!this->emitSetLocal(PT_Ptr, ResultOffset, E))
+      return false;
+  }
+
+  // Evaluate LHS and save value to LHSOffset.
+  if (!this->visit(LHS))
+    return false;
+  if (!this->emitSetLocal(PT_Ptr, LHSOffset, E))
+    return false;
+
+  // Same with RHS.
+  if (!this->visit(RHS))
+    return false;
+  if (!this->emitSetLocal(PT_Ptr, RHSOffset, E))
+    return false;
+
+  // Now we can get pointers to the LHS and RHS from the offsets above.
+  BinaryOperatorKind Op = E->getOpcode();
+  for (unsigned ElemIndex = 0; ElemIndex != 2; ++ElemIndex) {
+    // Result pointer for the store later.
+    if (!this->DiscardResult) {
+      if (!this->emitGetLocal(PT_Ptr, ResultOffset, E))
+        return false;
+    }
+
+    if (!this->emitGetLocal(PT_Ptr, LHSOffset, E))
+      return false;
+    if (!this->emitConstUint8(ElemIndex, E))
+      return false;
+    if (!this->emitArrayElemPtrPopUint8(E))
+      return false;
+    if (!this->emitLoadPop(LHSElemT, E))
+      return false;
+
+    if (!this->emitGetLocal(PT_Ptr, RHSOffset, E))
+      return false;
+    if (!this->emitConstUint8(ElemIndex, E))
+      return false;
+    if (!this->emitArrayElemPtrPopUint8(E))
+      return false;
+    if (!this->emitLoadPop(RHSElemT, E))
+      return false;
+
+    // The actual operation.
+    switch (Op) {
+    case BO_Add:
+      if (LHSElemT == PT_Float) {
+        if (!this->emitAddf(getRoundingMode(E), E))
+          return false;
+      } else {
+        if (!this->emitAdd(LHSElemT, E))
+          return false;
+      }
+      break;
+    case BO_Sub:
+      if (LHSElemT == PT_Float) {
+        if (!this->emitSubf(getRoundingMode(E), E))
+          return false;
+      } else {
+        if (!this->emitSub(LHSElemT, E))
+          return false;
+      }
+      break;
+
+    default:
+      return false;
+    }
+
+    if (!this->DiscardResult) {
+      // Initialize array element with the value we just computed.
+      if (!this->emitInitElemPop(LHSElemT, ElemIndex, E))
+        return false;
+    } else {
+      if (!this->emitPop(LHSElemT, E))
+        return false;
+    }
+  }
+  return true;
+}
+
 template <class Emitter>
 bool ByteCodeExprGen<Emitter>::VisitImplicitValueInitExpr(const ImplicitValueInitExpr *E) {
   QualType QT = E->getType();
@@ -671,6 +833,32 @@ bool ByteCodeExprGen<Emitter>::VisitInitListExpr(const InitListExpr *E) {
     return true;
   }
 
+  if (T->isAnyComplexType()) {
+    unsigned NumInits = E->getNumInits();
+    QualType ElemQT = E->getType()->getAs<ComplexType>()->getElementType();
+    PrimType ElemT = classifyPrim(ElemQT);
+    if (NumInits == 0) {
+      // Zero-initialize both elements.
+      for (unsigned I = 0; I < 2; ++I) {
+        if (!this->visitZeroInitializer(ElemT, ElemQT, E))
+          return false;
+        if (!this->emitInitElem(ElemT, I, E))
+          return false;
+      }
+    } else if (NumInits == 2) {
+      unsigned InitIndex = 0;
+      for (const Expr *Init : E->inits()) {
+        if (!this->visit(Init))
+          return false;
+
+        if (!this->emitInitElem(ElemT, InitIndex, E))
+          return false;
+        ++InitIndex;
+      }
+    }
+    return true;
+  }
+
   return false;
 }
 
@@ -1647,7 +1835,8 @@ template <class Emitter> bool ByteCodeExprGen<Emitter>::visit(const Expr *E) {
     return this->discard(E);
 
   // Create local variable to hold the return value.
-  if (!E->isGLValue() && !classify(E->getType())) {
+  if (!E->isGLValue() && !E->getType()->isAnyComplexType() &&
+      !classify(E->getType())) {
     std::optional<unsigned> LocalIndex = allocateLocal(E, /*IsExtended=*/true);
     if (!LocalIndex)
       return false;
@@ -1833,6 +2022,9 @@ bool ByteCodeExprGen<Emitter>::dereference(
     return Indirect(*T);
   }
 
+  if (LV->getType()->isAnyComplexType())
+    return visit(LV);
+
   return false;
 }
 
@@ -2092,15 +2284,33 @@ const Function *ByteCodeExprGen<Emitter>::getFunction(const FunctionDecl *FD) {
 template <class Emitter>
 bool ByteCodeExprGen<Emitter>::visitExpr(const Expr *E) {
   ExprScope<Emitter> RootScope(this);
+  // Void expressions.
+  if (E->getType()->isVoidType()) {
     if (!visit(E))
       return false;
-
-  if (E->getType()->isVoidType())
     return this->emitRetVoid(E);
+  }
 
-  if (std::optional<PrimType> T = classify(E))
+  // Expressions with a primitive return type.
+  if (std::optional<PrimType> T = classify(E)) {
+    if (!visit(E))
+      return false;
     return this->emitRet(*T, E);
+  }
+
+  // Expressions with a composite return type.
+  // For us, that means everything we don't
+  // have a PrimType for.
+  if (std::optional<unsigned> LocalOffset = this->allocateLocal(E)) {
+    if (!this->visitLocalInitializer(E, *LocalOffset))
+      return false;
+
+    if (!this->emitGetPtrLocal(*LocalOffset, E))
+      return false;
     return this->emitRetValue(E);
+  }
+
+  return false;
 }
 
 /// Toplevel visitDecl().
@@ -2550,8 +2760,36 @@ bool ByteCodeExprGen<Emitter>::VisitUnaryOperator(const UnaryOperator *E) {
     if (!this->visit(SubExpr))
       return false;
     return DiscardResult ? this->emitPop(*T, E) : this->emitComp(*T, E);
-  case UO_Real:   // __real x
-  case UO_Imag:   // __imag x
+  case UO_Real: { // __real x
+    assert(!T);
+    if (!this->visit(SubExpr))
+      return false;
+    if (!this->emitConstUint8(0, E))
+      return false;
+    if (!this->emitArrayElemPtrPopUint8(E))
+      return false;
+
+    // Since our _Complex implementation does not map to a primitive type,
+    // we sometimes have to do the lvalue-to-rvalue conversion here manually.
+    if (!SubExpr->isLValue())
+      return this->emitLoadPop(classifyPrim(E->getType()), E);
+    return true;
+  }
+  case UO_Imag: { // __imag x
+    assert(!T);
+    if (!this->visit(SubExpr))
+      return false;
+    if (!this->emitConstUint8(1, E))
+      return false;
+    if (!this->emitArrayElemPtrPopUint8(E))
+      return false;
+
+    // Since our _Complex implementation does not map to a primitive type,
+    // we sometimes have to do the lvalue-to-rvalue conversion here manually.
+    if (!SubExpr->isLValue())
+      return this->emitLoadPop(classifyPrim(E->getType()), E);
+    return true;
+  }
   case UO_Extension:
     return this->delegate(SubExpr);
   case UO_Coawait:

clang/lib/AST/Interp/ByteCodeExprGen.h

@@ -65,6 +65,7 @@ class ByteCodeExprGen : public ConstStmtVisitor<ByteCodeExprGen<Emitter>, bool>,
   bool VisitBinaryOperator(const BinaryOperator *E);
   bool VisitLogicalBinOp(const BinaryOperator *E);
   bool VisitPointerArithBinOp(const BinaryOperator *E);
+  bool VisitComplexBinOp(const BinaryOperator *E);
   bool VisitCXXDefaultArgExpr(const CXXDefaultArgExpr *E);
   bool VisitCallExpr(const CallExpr *E);
   bool VisitBuiltinCallExpr(const CallExpr *E);
@@ -285,6 +286,14 @@ class ByteCodeExprGen : public ConstStmtVisitor<ByteCodeExprGen<Emitter>, bool>,
   }
 
   bool emitPrimCast(PrimType FromT, PrimType ToT, QualType ToQT, const Expr *E);
+  std::optional<PrimType> classifyComplexElementType(QualType T) const {
+    assert(T->isAnyComplexType());
+
+    QualType ElemType = T->getAs<ComplexType>()->getElementType();
+
+    return this->classify(ElemType);
+  }
+
   bool emitRecordDestruction(const Descriptor *Desc);
   unsigned collectBaseOffset(const RecordType *BaseType,
                              const RecordType *DerivedType);

clang/lib/AST/Interp/Context.cpp

@@ -92,6 +92,9 @@ std::optional<PrimType> Context::classify(QualType T) const {
   if (T->isBooleanType())
     return PT_Bool;
 
+  if (T->isAnyComplexType())
+    return std::nullopt;
+
   if (T->isSignedIntegerOrEnumerationType()) {
     switch (Ctx.getIntWidth(T)) {
     case 64:
@@ -168,7 +171,7 @@ bool Context::Run(State &Parent, const Function *Func, APValue &Result) {
     }
 
     // State gets destroyed here, so the Stk.clear() below doesn't accidentally
-    // remove values the State's destructor might accedd.
+    // remove values the State's destructor might access.
   }
 
   Stk.clear();

clang/lib/AST/Interp/EvalEmitter.cpp

@@ -208,6 +208,27 @@ bool EvalEmitter::emitRetValue(const SourceInfo &Info) {
       }
       return Ok;
     }
+
+    // Complex types.
+    if (const auto *CT = Ty->getAs<ComplexType>()) {
+      QualType ElemTy = CT->getElementType();
+      std::optional<PrimType> ElemT = Ctx.classify(ElemTy);
+      assert(ElemT);
+
+      if (ElemTy->isIntegerType()) {
+        INT_TYPE_SWITCH(*ElemT, {
+          auto V1 = Ptr.atIndex(0).deref<T>();
+          auto V2 = Ptr.atIndex(1).deref<T>();
+          Result = APValue(V1.toAPSInt(), V2.toAPSInt());
+          return true;
+        });
+      } else if (ElemTy->isFloatingType()) {
+        Result = APValue(Ptr.atIndex(0).deref<Floating>().getAPFloat(),
+                         Ptr.atIndex(1).deref<Floating>().getAPFloat());
+        return true;
+      }
+      return false;
+    }
     llvm_unreachable("invalid value to return");
   };
 

clang/lib/AST/Interp/IntegralAP.h

@@ -182,17 +182,13 @@ template <bool Signed> class IntegralAP final {
   }
 
   static bool increment(IntegralAP A, IntegralAP *R) {
-    // FIXME: Implement.
-    assert(false);
-    *R = IntegralAP(A.V - 1);
-    return false;
+    IntegralAP<Signed> One(1, A.bitWidth());
+    return add(A, One, A.bitWidth() + 1, R);
   }
 
   static bool decrement(IntegralAP A, IntegralAP *R) {
-    // FIXME: Implement.
-    assert(false);
-    *R = IntegralAP(A.V - 1);
-    return false;
+    IntegralAP<Signed> One(1, A.bitWidth());
+    return sub(A, One, A.bitWidth() + 1, R);
   }
 
   static bool add(IntegralAP A, IntegralAP B, unsigned OpBits, IntegralAP *R) {

clang/lib/AST/Interp/Interp.cpp

@@ -350,11 +350,6 @@ bool CheckCallable(InterpState &S, CodePtr OpPC, const Function *F) {
   }
 
   if (!F->isConstexpr()) {
-    // Don't emit anything if we're checking for a potential constant
-    // expression. That will happen later when actually executing.
-    if (S.checkingPotentialConstantExpression())
-      return false;
-
     const SourceLocation &Loc = S.Current->getLocation(OpPC);
     if (S.getLangOpts().CPlusPlus11) {
       const FunctionDecl *DiagDecl = F->getDecl();
@@ -371,13 +366,21 @@ bool CheckCallable(InterpState &S, CodePtr OpPC, const Function *F) {
       // FIXME: If DiagDecl is an implicitly-declared special member function
       // or an inheriting constructor, we should be much more explicit about why
       // it's not constexpr.
-      if (CD && CD->isInheritingConstructor())
+      if (CD && CD->isInheritingConstructor()) {
         S.FFDiag(Loc, diag::note_constexpr_invalid_inhctor, 1)
           << CD->getInheritedConstructor().getConstructor()->getParent();
-      else
+        S.Note(DiagDecl->getLocation(), diag::note_declared_at);
+      } else {
+        // Don't emit anything if the function isn't defined and we're checking
+        // for a constant expression. It might be defined at the point we're
+        // actually calling it.
+        if (!DiagDecl->isDefined() && S.checkingPotentialConstantExpression())
+          return false;
+
         S.FFDiag(Loc, diag::note_constexpr_invalid_function, 1)
           << DiagDecl->isConstexpr() << (bool)CD << DiagDecl;
         S.Note(DiagDecl->getLocation(), diag::note_declared_at);
+      }
     } else {
       S.FFDiag(Loc, diag::note_invalid_subexpr_in_const_expr);
     }

clang/lib/AST/Interp/Interp.h

@@ -1619,7 +1619,11 @@ bool CastFloatingIntegral(InterpState &S, CodePtr OpPC) {
       QualType Type = E->getType();
 
       S.CCEDiag(E, diag::note_constexpr_overflow) << F.getAPFloat() << Type;
-      return S.noteUndefinedBehavior();
+      if (S.noteUndefinedBehavior()) {
+        S.Stk.push<T>(T(Result));
+        return true;
+      }
+      return false;
     }
 
     S.Stk.push<T>(T(Result));
@@ -1822,10 +1826,12 @@ inline bool ArrayElemPtr(InterpState &S, CodePtr OpPC) {
 /// Just takes a pointer and checks if its' an incomplete
 /// array type.
 inline bool ArrayDecay(InterpState &S, CodePtr OpPC) {
-  const Pointer &Ptr = S.Stk.peek<Pointer>();
+  const Pointer &Ptr = S.Stk.pop<Pointer>();
 
-  if (!Ptr.isUnknownSizeArray())
+  if (!Ptr.isUnknownSizeArray()) {
+    S.Stk.push<Pointer>(Ptr.atIndex(0));
     return true;
+  }
 
   const SourceInfo &E = S.Current->getSource(OpPC);
   S.FFDiag(E, diag::note_constexpr_unsupported_unsized_array);

clang/lib/AST/Interp/InterpBuiltin.cpp

@@ -579,6 +579,40 @@ static bool interp__builtin_expect(InterpState &S, CodePtr OpPC,
   return true;
 }
 
+/// rotateleft(value, amount)
+static bool interp__builtin_rotate(InterpState &S, CodePtr OpPC,
+                                   const InterpFrame *Frame,
+                                   const Function *Func, const CallExpr *Call,
+                                   bool Right) {
+  PrimType ArgT = *S.getContext().classify(Call->getArg(0)->getType());
+  assert(ArgT == *S.getContext().classify(Call->getArg(1)->getType()));
+
+  APSInt Amount = peekToAPSInt(S.Stk, ArgT);
+  APSInt Value = peekToAPSInt(S.Stk, ArgT, align(primSize(ArgT)) * 2);
+
+  APSInt Result;
+  if (Right)
+    Result = APSInt(Value.rotr(Amount.urem(Value.getBitWidth())),
+                    /*IsUnsigned=*/true);
+  else // Left.
+    Result = APSInt(Value.rotl(Amount.urem(Value.getBitWidth())),
+                    /*IsUnsigned=*/true);
+
+  pushAPSInt(S, Result);
+  return true;
+}
+
+static bool interp__builtin_ffs(InterpState &S, CodePtr OpPC,
+                                const InterpFrame *Frame, const Function *Func,
+                                const CallExpr *Call) {
+  PrimType ArgT = *S.getContext().classify(Call->getArg(0)->getType());
+  APSInt Value = peekToAPSInt(S.Stk, ArgT);
+
+  uint64_t N = Value.countr_zero();
+  pushInt(S, N == Value.getBitWidth() ? 0 : N + 1);
+  return true;
+}
+
 bool InterpretBuiltin(InterpState &S, CodePtr OpPC, const Function *F,
                       const CallExpr *Call) {
   InterpFrame *Frame = S.Current;
@@ -754,6 +788,39 @@ bool InterpretBuiltin(InterpState &S, CodePtr OpPC, const Function *F,
       return false;
     break;
 
+  case Builtin::BI__builtin_rotateleft8:
+  case Builtin::BI__builtin_rotateleft16:
+  case Builtin::BI__builtin_rotateleft32:
+  case Builtin::BI__builtin_rotateleft64:
+  case Builtin::BI_rotl8: // Microsoft variants of rotate left
+  case Builtin::BI_rotl16:
+  case Builtin::BI_rotl:
+  case Builtin::BI_lrotl:
+  case Builtin::BI_rotl64:
+    if (!interp__builtin_rotate(S, OpPC, Frame, F, Call, /*Right=*/false))
+      return false;
+    break;
+
+  case Builtin::BI__builtin_rotateright8:
+  case Builtin::BI__builtin_rotateright16:
+  case Builtin::BI__builtin_rotateright32:
+  case Builtin::BI__builtin_rotateright64:
+  case Builtin::BI_rotr8: // Microsoft variants of rotate right
+  case Builtin::BI_rotr16:
+  case Builtin::BI_rotr:
+  case Builtin::BI_lrotr:
+  case Builtin::BI_rotr64:
+    if (!interp__builtin_rotate(S, OpPC, Frame, F, Call, /*Right=*/true))
+      return false;
+    break;
+
+  case Builtin::BI__builtin_ffs:
+  case Builtin::BI__builtin_ffsl:
+  case Builtin::BI__builtin_ffsll:
+    if (!interp__builtin_ffs(S, OpPC, Frame, F, Call))
+      return false;
+    break;
+
   default:
     return false;
   }

clang/lib/AST/Interp/InterpFrame.cpp

@@ -228,7 +228,7 @@ Pointer InterpFrame::getParamPointer(unsigned Off) {
 SourceInfo InterpFrame::getSource(CodePtr PC) const {
   // Implicitly created functions don't have any code we could point at,
   // so return the call site.
-  if (Func && Func->getDecl()->isImplicit() && Caller)
+  if (Func && (!Func->hasBody() || Func->getDecl()->isImplicit()) && Caller)
     return Caller->getSource(RetPC);
 
   return S.getSource(Func, PC);
@@ -243,7 +243,7 @@ SourceLocation InterpFrame::getLocation(CodePtr PC) const {
 }
 
 SourceRange InterpFrame::getRange(CodePtr PC) const {
-  if (Func && Func->getDecl()->isImplicit() && Caller)
+  if (Func && (!Func->hasBody() || Func->getDecl()->isImplicit()) && Caller)
     return Caller->getRange(RetPC);
 
   return S.getRange(Func, PC);

clang/lib/AST/Mangle.cpp

@@ -147,7 +147,7 @@ void MangleContext::mangleName(GlobalDecl GD, raw_ostream &Out) {
 
     // If the label isn't literal, or if this is an alias for an LLVM intrinsic,
     // do not add a "\01" prefix.
-    if (!ALA->getIsLiteralLabel() || ALA->getLabel().startswith("llvm.")) {
+    if (!ALA->getIsLiteralLabel() || ALA->getLabel().starts_with("llvm.")) {
       Out << ALA->getLabel();
       return;
     }

clang/lib/AST/MicrosoftMangle.cpp

@@ -63,7 +63,7 @@ struct msvc_hashing_ostream : public llvm::raw_svector_ostream {
       : llvm::raw_svector_ostream(Buffer), OS(OS) {}
   ~msvc_hashing_ostream() override {
     StringRef MangledName = str();
-    bool StartsWithEscape = MangledName.startswith("\01");
+    bool StartsWithEscape = MangledName.starts_with("\01");
     if (StartsWithEscape)
       MangledName = MangledName.drop_front(1);
     if (MangledName.size() < 4096) {
@@ -3809,14 +3809,14 @@ void MicrosoftMangleContextImpl::mangleCXXRTTICompleteObjectLocator(
   llvm::raw_svector_ostream Stream(VFTableMangling);
   mangleCXXVFTable(Derived, BasePath, Stream);
 
-  if (VFTableMangling.startswith("??@")) {
-    assert(VFTableMangling.endswith("@"));
+  if (VFTableMangling.starts_with("??@")) {
+    assert(VFTableMangling.ends_with("@"));
     Out << VFTableMangling << "??_R4@";
     return;
   }
 
-  assert(VFTableMangling.startswith("??_7") ||
-         VFTableMangling.startswith("??_S"));
+  assert(VFTableMangling.starts_with("??_7") ||
+         VFTableMangling.starts_with("??_S"));
 
   Out << "??_R4" << VFTableMangling.str().drop_front(4);
 }

clang/lib/AST/PrintfFormatString.cpp

@@ -140,7 +140,7 @@ static PrintfSpecifierResult ParsePrintfSpecifier(FormatStringHandler &H,
         // Set the privacy flag if the privacy annotation in the
         // comma-delimited segment is at least as strict as the privacy
         // annotations in previous comma-delimited segments.
-        if (MatchedStr.startswith("mask")) {
+        if (MatchedStr.starts_with("mask")) {
           StringRef MaskType = MatchedStr.substr(sizeof("mask.") - 1);
           unsigned Size = MaskType.size();
           if (Warn && (Size == 0 || Size > 8))

clang/lib/AST/RawCommentList.cpp

@@ -141,8 +141,8 @@ RawComment::RawComment(const SourceManager &SourceMgr, SourceRange SR,
     Kind = K.first;
     IsTrailingComment |= K.second;
 
-    IsAlmostTrailingComment = RawText.startswith("//<") ||
-                                 RawText.startswith("/*<");
+    IsAlmostTrailingComment =
+        RawText.starts_with("//<") || RawText.starts_with("/*<");
   } else {
     Kind = RCK_Merged;
     IsTrailingComment =

clang/lib/AST/RecordLayoutBuilder.cpp

@@ -2942,8 +2942,8 @@ void MicrosoftRecordLayoutBuilder::layoutNonVirtualBase(
   }
 
   if (!FoundBase) {
-    if (MDCUsesEBO && BaseDecl->isEmpty()) {
-      assert(BaseLayout.getNonVirtualSize() == CharUnits::Zero());
+    if (MDCUsesEBO && BaseDecl->isEmpty() &&
+        (BaseLayout.getNonVirtualSize() == CharUnits::Zero())) {
       BaseOffset = CharUnits::Zero();
     } else {
       // Otherwise, lay the base out at the end of the MDC.

clang/lib/AST/Stmt.cpp

@@ -811,11 +811,12 @@ std::string MSAsmStmt::generateAsmString(const ASTContext &C) const {
     StringRef Instruction = Pieces[I];
     // For vex/vex2/vex3/evex masm style prefix, convert it to att style
     // since we don't support masm style prefix in backend.
-    if (Instruction.startswith("vex "))
+    if (Instruction.starts_with("vex "))
       MSAsmString += '{' + Instruction.substr(0, 3).str() + '}' +
                      Instruction.substr(3).str();
-    else if (Instruction.startswith("vex2 ") ||
-             Instruction.startswith("vex3 ") || Instruction.startswith("evex "))
+    else if (Instruction.starts_with("vex2 ") ||
+             Instruction.starts_with("vex3 ") ||
+             Instruction.starts_with("evex "))
       MSAsmString += '{' + Instruction.substr(0, 4).str() + '}' +
                      Instruction.substr(4).str();
     else

clang/lib/ASTMatchers/ASTMatchersInternal.cpp

@@ -480,11 +480,11 @@ HasNameMatcher::HasNameMatcher(std::vector<std::string> N)
 
 static bool consumeNameSuffix(StringRef &FullName, StringRef Suffix) {
   StringRef Name = FullName;
-  if (!Name.endswith(Suffix))
+  if (!Name.ends_with(Suffix))
     return false;
   Name = Name.drop_back(Suffix.size());
   if (!Name.empty()) {
-    if (!Name.endswith("::"))
+    if (!Name.ends_with("::"))
       return false;
     Name = Name.drop_back(2);
   }
@@ -530,7 +530,7 @@ class PatternSet {
   PatternSet(ArrayRef<std::string> Names) {
     Patterns.reserve(Names.size());
     for (StringRef Name : Names)
-      Patterns.push_back({Name, Name.startswith("::")});
+      Patterns.push_back({Name, Name.starts_with("::")});
   }
 
   /// Consumes the name suffix from each pattern in the set and removes the ones
@@ -652,11 +652,11 @@ bool HasNameMatcher::matchesNodeFullSlow(const NamedDecl &Node) const {
     const StringRef FullName = OS.str();
 
     for (const StringRef Pattern : Names) {
-      if (Pattern.startswith("::")) {
+      if (Pattern.starts_with("::")) {
         if (FullName == Pattern)
           return true;
-      } else if (FullName.endswith(Pattern) &&
-                 FullName.drop_back(Pattern.size()).endswith("::")) {
+      } else if (FullName.ends_with(Pattern) &&
+                 FullName.drop_back(Pattern.size()).ends_with("::")) {
         return true;
       }
     }

clang/lib/ASTMatchers/Dynamic/Parser.cpp

@@ -187,10 +187,10 @@ class Parser::CodeTokenizer {
             break;
           ++TokenLength;
         }
-        if (TokenLength == 4 && Code.startswith("true")) {
+        if (TokenLength == 4 && Code.starts_with("true")) {
           Result.Kind = TokenInfo::TK_Literal;
           Result.Value = true;
-        } else if (TokenLength == 5 && Code.startswith("false")) {
+        } else if (TokenLength == 5 && Code.starts_with("false")) {
           Result.Kind = TokenInfo::TK_Literal;
           Result.Value = false;
         } else {
@@ -737,7 +737,7 @@ bool Parser::parseMatcherExpressionImpl(const TokenInfo &NameToken,
 // Completions minus the prefix.
 void Parser::addCompletion(const TokenInfo &CompToken,
                            const MatcherCompletion& Completion) {
-  if (StringRef(Completion.TypedText).startswith(CompToken.Text) &&
+  if (StringRef(Completion.TypedText).starts_with(CompToken.Text) &&
       Completion.Specificity > 0) {
     Completions.emplace_back(Completion.TypedText.substr(CompToken.Text.size()),
                              Completion.MatcherDecl, Completion.Specificity);

clang/lib/Analysis/BodyFarm.cpp

@@ -726,8 +726,8 @@ Stmt *BodyFarm::getBody(const FunctionDecl *D) {
       FF = nullptr;
       break;
     }
-  } else if (Name.startswith("OSAtomicCompareAndSwap") ||
-             Name.startswith("objc_atomicCompareAndSwap")) {
+  } else if (Name.starts_with("OSAtomicCompareAndSwap") ||
+             Name.starts_with("objc_atomicCompareAndSwap")) {
     FF = create_OSAtomicCompareAndSwap;
   } else if (Name == "call_once" && D->getDeclContext()->isStdNamespace()) {
     FF = create_call_once;

clang/lib/Analysis/CallGraph.cpp

@@ -168,7 +168,7 @@ bool CallGraph::includeCalleeInGraph(const Decl *D) {
       return false;
 
     IdentifierInfo *II = FD->getIdentifier();
-    if (II && II->getName().startswith("__inline"))
+    if (II && II->getName().starts_with("__inline"))
       return false;
   }
 

clang/lib/Analysis/CalledOnceCheck.cpp

@@ -973,7 +973,7 @@ class CalledOnceChecker : public ConstStmtVisitor<CalledOnceChecker> {
   /// Return true if the given name has conventional suffixes.
   static bool hasConventionalSuffix(llvm::StringRef Name) {
     return llvm::any_of(CONVENTIONAL_SUFFIXES, [Name](llvm::StringRef Suffix) {
-      return Name.endswith(Suffix);
+      return Name.ends_with(Suffix);
     });
   }
 

clang/lib/Analysis/CocoaConventions.cpp

@@ -26,10 +26,10 @@ bool cocoa::isRefType(QualType RetTy, StringRef Prefix,
   // Recursively walk the typedef stack, allowing typedefs of reference types.
   while (const TypedefType *TD = RetTy->getAs<TypedefType>()) {
     StringRef TDName = TD->getDecl()->getIdentifier()->getName();
-    if (TDName.startswith(Prefix) && TDName.endswith("Ref"))
+    if (TDName.starts_with(Prefix) && TDName.ends_with("Ref"))
       return true;
     // XPC unfortunately uses CF-style function names, but aren't CF types.
-    if (TDName.startswith("xpc_"))
+    if (TDName.starts_with("xpc_"))
       return false;
     RetTy = TD->getDecl()->getUnderlyingType();
   }
@@ -43,7 +43,7 @@ bool cocoa::isRefType(QualType RetTy, StringRef Prefix,
     return false;
 
   // Does the name start with the prefix?
-  return Name.startswith(Prefix);
+  return Name.starts_with(Prefix);
 }
 
 /// Returns true when the passed-in type is a CF-style reference-counted
@@ -127,10 +127,9 @@ bool coreFoundation::followsCreateRule(const FunctionDecl *fn) {
     // Scan for *lowercase* 'reate' or 'opy', followed by no lowercase
     // character.
     StringRef suffix = functionName.substr(it - start);
-    if (suffix.startswith("reate")) {
+    if (suffix.starts_with("reate")) {
       it += 5;
-    }
-    else if (suffix.startswith("opy")) {
+    } else if (suffix.starts_with("opy")) {
       it += 3;
     } else {
       // Keep scanning.

clang/lib/Analysis/FlowSensitive/Models/ChromiumCheckModel.cpp

@@ -43,7 +43,7 @@ bool isCheckLikeMethod(llvm::SmallDenseSet<const CXXMethodDecl *> &CheckDecls,
       return false;
 
     for (const CXXMethodDecl *M : ParentClass->methods())
-      if (M->getDeclName().isIdentifier() && M->getName().endswith("Check"))
+      if (M->getDeclName().isIdentifier() && M->getName().ends_with("Check"))
         CheckDecls.insert(M);
   }
 

clang/lib/Analysis/RetainSummaryManager.cpp

@@ -174,7 +174,7 @@ static bool isOSObjectPtr(QualType QT) {
 }
 
 static bool isISLObjectRef(QualType Ty) {
-  return StringRef(Ty.getAsString()).startswith("isl_");
+  return StringRef(Ty.getAsString()).starts_with("isl_");
 }
 
 static bool isOSIteratorSubclass(const Decl *D) {
@@ -255,13 +255,13 @@ RetainSummaryManager::getSummaryForOSObject(const FunctionDecl *FD,
       // TODO: Add support for the slightly common *Matching(table) idiom.
       // Cf. IOService::nameMatching() etc. - these function have an unusual
       // contract of returning at +0 or +1 depending on their last argument.
-      if (FName.endswith("Matching")) {
+      if (FName.ends_with("Matching")) {
         return getPersistentStopSummary();
       }
 
       // All objects returned with functions *not* starting with 'get',
       // or iterators, are returned at +1.
-      if ((!FName.startswith("get") && !FName.startswith("Get")) ||
+      if ((!FName.starts_with("get") && !FName.starts_with("Get")) ||
           isOSIteratorSubclass(PD)) {
         return getOSSummaryCreateRule(FD);
       } else {
@@ -392,9 +392,9 @@ const RetainSummary *RetainSummaryManager::getSummaryForObjCOrCFObject(
     return getPersistentSummary(RetEffect::MakeNoRet(),
                                 ScratchArgs,
                                 ArgEffect(DoNothing), ArgEffect(DoNothing));
-  } else if (FName.startswith("NSLog")) {
+  } else if (FName.starts_with("NSLog")) {
     return getDoNothingSummary();
-  } else if (FName.startswith("NS") && FName.contains("Insert")) {
+  } else if (FName.starts_with("NS") && FName.contains("Insert")) {
     // Allowlist NSXXInsertXX, for example NSMapInsertIfAbsent, since they can
     // be deallocated by NSMapRemove.
     ScratchArgs = AF.add(ScratchArgs, 1, ArgEffect(StopTracking));
@@ -453,9 +453,9 @@ const RetainSummary *RetainSummaryManager::getSummaryForObjCOrCFObject(
 
   // Check for release functions, the only kind of functions that we care
   // about that don't return a pointer type.
-  if (FName.startswith("CG") || FName.startswith("CF")) {
+  if (FName.starts_with("CG") || FName.starts_with("CF")) {
     // Test for 'CGCF'.
-    FName = FName.substr(FName.startswith("CGCF") ? 4 : 2);
+    FName = FName.substr(FName.starts_with("CGCF") ? 4 : 2);
 
     if (isRelease(FD, FName))
       return getUnarySummary(FT, DecRef);

clang/lib/Analysis/UnsafeBufferUsage.cpp

@@ -1028,6 +1028,46 @@ class UPCPreIncrementGadget : public FixableGadget {
   }
 };
 
+// Representing a pointer type expression of the form `Ptr += n` in an
+// Unspecified Untyped Context (UUC):
+class UUCAddAssignGadget : public FixableGadget {
+private:
+  static constexpr const char *const UUCAddAssignTag =
+      "PointerAddAssignUnderUUC";
+  static constexpr const char *const OffsetTag = "Offset";
+
+  const BinaryOperator *Node; // the `Ptr += n` node
+  const Expr *Offset = nullptr;
+
+public:
+  UUCAddAssignGadget(const MatchFinder::MatchResult &Result)
+      : FixableGadget(Kind::UUCAddAssign),
+        Node(Result.Nodes.getNodeAs<BinaryOperator>(UUCAddAssignTag)),
+        Offset(Result.Nodes.getNodeAs<Expr>(OffsetTag)) {
+    assert(Node != nullptr && "Expecting a non-null matching result");
+  }
+
+  static bool classof(const Gadget *G) {
+    return G->getKind() == Kind::UUCAddAssign;
+  }
+
+  static Matcher matcher() {
+    return stmt(isInUnspecifiedUntypedContext(expr(ignoringImpCasts(
+        binaryOperator(hasOperatorName("+="),
+                       hasLHS(declRefExpr(toSupportedVariable())),
+                       hasRHS(expr().bind(OffsetTag)))
+            .bind(UUCAddAssignTag)))));
+  }
+
+  virtual std::optional<FixItList> getFixits(const Strategy &S) const override;
+
+  virtual const Stmt *getBaseStmt() const override { return Node; }
+
+  virtual DeclUseList getClaimedVarUseSites() const override {
+    return {dyn_cast<DeclRefExpr>(Node->getLHS())};
+  }
+};
+
 // Representing a fixable expression of the form `*(ptr + 123)` or `*(123 +
 // ptr)`:
 class DerefSimplePtrArithFixableGadget : public FixableGadget {
@@ -1312,6 +1352,16 @@ PointerInitGadget::getFixits(const Strategy &S) const {
   return std::nullopt;
 }
 
+static bool isNonNegativeIntegerExpr(const Expr *Expr, const VarDecl *VD,
+                                     const ASTContext &Ctx) {
+  if (auto ConstVal = Expr->getIntegerConstantExpr(Ctx)) {
+    if (ConstVal->isNegative())
+      return false;
+  } else if (!Expr->getType()->isUnsignedIntegerType())
+    return false;
+  return true;
+}
+
 std::optional<FixItList>
 ULCArraySubscriptGadget::getFixits(const Strategy &S) const {
   if (const auto *DRE =
@@ -1319,14 +1369,12 @@ ULCArraySubscriptGadget::getFixits(const Strategy &S) const {
     if (const auto *VD = dyn_cast<VarDecl>(DRE->getDecl())) {
       switch (S.lookup(VD)) {
       case Strategy::Kind::Span: {
+
         // If the index has a negative constant value, we give up as no valid
         // fix-it can be generated:
         const ASTContext &Ctx = // FIXME: we need ASTContext to be passed in!
             VD->getASTContext();
-        if (auto ConstVal = Node->getIdx()->getIntegerConstantExpr(Ctx)) {
-          if (ConstVal->isNegative())
-            return std::nullopt;
-        } else if (!Node->getIdx()->getType()->isUnsignedIntegerType())
+        if (!isNonNegativeIntegerExpr(Node->getIdx(), VD, Ctx))
           return std::nullopt;
         // no-op is a good fix-it, otherwise
         return FixItList{};
@@ -1405,10 +1453,8 @@ static std::optional<SourceLocation> getPastLoc(const NodeTy *Node,
                                                 const LangOptions &LangOpts) {
   SourceLocation Loc =
       Lexer::getLocForEndOfToken(Node->getEndLoc(), 0, SM, LangOpts);
-
   if (Loc.isValid())
     return Loc;
-
   return std::nullopt;
 }
 
@@ -1488,7 +1534,7 @@ static bool hasUnsupportedSpecifiers(const VarDecl *VD,
 // returned by this function is the last location of the last token.
 static SourceRange getSourceRangeToTokenEnd(const Decl *D,
                                             const SourceManager &SM,
-                                            LangOptions LangOpts) {
+                                            const LangOptions &LangOpts) {
   SourceLocation Begin = D->getBeginLoc();
   SourceLocation
     End = // `D->getEndLoc` should always return the starting location of the
@@ -1766,6 +1812,47 @@ fixUPCAddressofArraySubscriptWithSpan(const UnaryOperator *Node) {
       FixItHint::CreateReplacement(Node->getSourceRange(), SS.str())};
 }
 
+std::optional<FixItList>
+UUCAddAssignGadget::getFixits(const Strategy &S) const {
+  DeclUseList DREs = getClaimedVarUseSites();
+
+  if (DREs.size() != 1)
+    return std::nullopt; // In cases of `Ptr += n` where `Ptr` is not a DRE, we
+                         // give up
+  if (const VarDecl *VD = dyn_cast<VarDecl>(DREs.front()->getDecl())) {
+    if (S.lookup(VD) == Strategy::Kind::Span) {
+      FixItList Fixes;
+
+      const Stmt *AddAssignNode = getBaseStmt();
+      StringRef varName = VD->getName();
+      const ASTContext &Ctx = VD->getASTContext();
+
+      if (!isNonNegativeIntegerExpr(Offset, VD, Ctx))
+        return std::nullopt;
+
+      // To transform UUC(p += n) to UUC(p = p.subspan(..)):
+      bool NotParenExpr =
+          (Offset->IgnoreParens()->getBeginLoc() == Offset->getBeginLoc());
+      std::string SS = varName.str() + " = " + varName.str() + ".subspan";
+      if (NotParenExpr)
+        SS += "(";
+
+      std::optional<SourceLocation> AddAssignLocation = getEndCharLoc(
+          AddAssignNode, Ctx.getSourceManager(), Ctx.getLangOpts());
+      if (!AddAssignLocation)
+        return std::nullopt;
+
+      Fixes.push_back(FixItHint::CreateReplacement(
+          SourceRange(AddAssignNode->getBeginLoc(), Node->getOperatorLoc()),
+          SS));
+      if (NotParenExpr)
+        Fixes.push_back(FixItHint::CreateInsertion(
+            Offset->getEndLoc().getLocWithOffset(1), ")"));
+      return Fixes;
+    }
+  }
+  return std::nullopt; // Not in the cases that we can handle for now, give up.
+}
 
 std::optional<FixItList> UPCPreIncrementGadget::getFixits(const Strategy &S) const {
   DeclUseList DREs = getClaimedVarUseSites();

clang/lib/Basic/Attributes.cpp

@@ -33,7 +33,7 @@ int clang::hasAttribute(AttributeCommonInfo::Syntax Syntax,
                         const TargetInfo &Target, const LangOptions &LangOpts) {
   StringRef Name = Attr->getName();
   // Normalize the attribute name, __foo__ becomes foo.
-  if (Name.size() >= 4 && Name.startswith("__") && Name.endswith("__"))
+  if (Name.size() >= 4 && Name.starts_with("__") && Name.ends_with("__"))
     Name = Name.substr(2, Name.size() - 4);
 
   // Normalize the scope name, but only for gnu and clang attributes.
@@ -103,8 +103,8 @@ static StringRef normalizeAttrName(const IdentifierInfo *Name,
        (NormalizedScopeName.empty() || NormalizedScopeName == "gnu" ||
         NormalizedScopeName == "clang"));
   StringRef AttrName = Name->getName();
-  if (ShouldNormalize && AttrName.size() >= 4 && AttrName.startswith("__") &&
-      AttrName.endswith("__"))
+  if (ShouldNormalize && AttrName.size() >= 4 && AttrName.starts_with("__") &&
+      AttrName.ends_with("__"))
     AttrName = AttrName.slice(2, AttrName.size() - 2);
 
   return AttrName;

clang/lib/Basic/Cuda.cpp

@@ -39,6 +39,8 @@ static const CudaVersionMapEntry CudaNameVersionMap[] = {
     CUDA_ENTRY(11, 8),
     CUDA_ENTRY(12, 0),
     CUDA_ENTRY(12, 1),
+    CUDA_ENTRY(12, 2),
+    CUDA_ENTRY(12, 3),
     {"", CudaVersion::NEW, llvm::VersionTuple(std::numeric_limits<int>::max())},
     {"unknown", CudaVersion::UNKNOWN, {}} // End of list tombstone.
 };
@@ -93,6 +95,7 @@ static const CudaArchToStringMap arch_names[] = {
     SM(87),                          // Jetson/Drive AGX Orin
     SM(89),                          // Ada Lovelace
     SM(90),                          // Hopper
+    SM(90a),                         // Hopper
     GFX(600),  // gfx600
     GFX(601),  // gfx601
     GFX(602),  // gfx602
@@ -209,6 +212,8 @@ CudaVersion MinVersionForCudaArch(CudaArch A) {
   case CudaArch::SM_89:
   case CudaArch::SM_90:
     return CudaVersion::CUDA_118;
+  case CudaArch::SM_90a:
+    return CudaVersion::CUDA_120;
   default:
     llvm_unreachable("invalid enum");
   }

clang/lib/Basic/DiagnosticIDs.cpp

@@ -853,5 +853,5 @@ bool DiagnosticIDs::isUnrecoverable(unsigned DiagID) const {
 
 bool DiagnosticIDs::isARCDiagnostic(unsigned DiagID) {
   unsigned cat = getCategoryNumberForDiag(DiagID);
-  return DiagnosticIDs::getCategoryNameFromID(cat).startswith("ARC ");
+  return DiagnosticIDs::getCategoryNameFromID(cat).starts_with("ARC ");
 }

clang/lib/Basic/IdentifierTable.cpp

@@ -604,7 +604,7 @@ LLVM_DUMP_METHOD void Selector::dump() const { print(llvm::errs()); }
 static bool startsWithWord(StringRef name, StringRef word) {
   if (name.size() < word.size()) return false;
   return ((name.size() == word.size() || !isLowercase(name[word.size()])) &&
-          name.startswith(word));
+          name.starts_with(word));
 }
 
 ObjCMethodFamily Selector::getMethodFamilyImpl(Selector sel) {
@@ -742,7 +742,7 @@ SelectorTable::constructSetterSelector(IdentifierTable &Idents,
 
 std::string SelectorTable::getPropertyNameFromSetterSelector(Selector Sel) {
   StringRef Name = Sel.getNameForSlot(0);
-  assert(Name.startswith("set") && "invalid setter name");
+  assert(Name.starts_with("set") && "invalid setter name");
   return (Twine(toLowercase(Name[3])) + Name.drop_front(4)).str();
 }
 

clang/lib/Basic/Module.cpp

@@ -89,7 +89,7 @@ static bool isPlatformEnvironment(const TargetInfo &Target, StringRef Feature) {
   // where both are valid examples of the same platform+environment but in the
   // variant (2) the simulator is hardcoded as part of the platform name. Both
   // forms above should match for "iossimulator" requirement.
-  if (Target.getTriple().isOSDarwin() && PlatformEnv.endswith("simulator"))
+  if (Target.getTriple().isOSDarwin() && PlatformEnv.ends_with("simulator"))
     return PlatformEnv == Feature || CmpPlatformEnv(PlatformEnv, Feature);
 
   return PlatformEnv == Feature;
@@ -166,7 +166,8 @@ bool Module::isForBuilding(const LangOptions &LangOpts) const {
   // for either.
   if (!LangOpts.isCompilingModule() && getTopLevelModule()->IsFramework &&
       CurrentModule == LangOpts.ModuleName &&
-      !CurrentModule.endswith("_Private") && TopLevelName.endswith("_Private"))
+      !CurrentModule.ends_with("_Private") &&
+      TopLevelName.ends_with("_Private"))
     TopLevelName = TopLevelName.drop_back(8);
 
   return TopLevelName == CurrentModule;

clang/lib/Basic/Sarif.cpp

@@ -74,7 +74,7 @@ static std::string fileNameToURI(StringRef Filename) {
 
   // Get the root name to see if it has a URI authority.
   StringRef Root = sys::path::root_name(Filename);
-  if (Root.startswith("//")) {
+  if (Root.starts_with("//")) {
     // There is an authority, so add it to the URI.
     Ret += Root.drop_front(2).str();
   } else if (!Root.empty()) {

clang/lib/Basic/TargetInfo.cpp

@@ -551,26 +551,26 @@ ParsedTargetAttr TargetInfo::parseTargetAttr(StringRef Features) const {
     // TODO: Support the fpmath option. It will require checking
     // overall feature validity for the function with the rest of the
     // attributes on the function.
-    if (Feature.startswith("fpmath="))
+    if (Feature.starts_with("fpmath="))
       continue;
 
-    if (Feature.startswith("branch-protection=")) {
+    if (Feature.starts_with("branch-protection=")) {
       Ret.BranchProtection = Feature.split('=').second.trim();
       continue;
     }
 
     // While we're here iterating check for a different target cpu.
-    if (Feature.startswith("arch=")) {
+    if (Feature.starts_with("arch=")) {
       if (!Ret.CPU.empty())
         Ret.Duplicate = "arch=";
       else
         Ret.CPU = Feature.split("=").second.trim();
-    } else if (Feature.startswith("tune=")) {
+    } else if (Feature.starts_with("tune=")) {
       if (!Ret.Tune.empty())
         Ret.Duplicate = "tune=";
       else
         Ret.Tune = Feature.split("=").second.trim();
-    } else if (Feature.startswith("no-"))
+    } else if (Feature.starts_with("no-"))
       Ret.Features.push_back("-" + Feature.split("-").second.str());
     else
       Ret.Features.push_back("+" + Feature.str());

clang/lib/Basic/Targets/AArch64.cpp

@@ -574,11 +574,12 @@ void AArch64TargetInfo::getTargetDefines(const LangOptions &Opts,
   else if (*ArchInfo == llvm::AArch64::ARMV9_5A)
     getTargetDefinesARMV95A(Opts, Builder);
 
-  // All of the __sync_(bool|val)_compare_and_swap_(1|2|4|8) builtins work.
+  // All of the __sync_(bool|val)_compare_and_swap_(1|2|4|8|16) builtins work.
   Builder.defineMacro("__GCC_HAVE_SYNC_COMPARE_AND_SWAP_1");
   Builder.defineMacro("__GCC_HAVE_SYNC_COMPARE_AND_SWAP_2");
   Builder.defineMacro("__GCC_HAVE_SYNC_COMPARE_AND_SWAP_4");
   Builder.defineMacro("__GCC_HAVE_SYNC_COMPARE_AND_SWAP_8");
+  Builder.defineMacro("__GCC_HAVE_SYNC_COMPARE_AND_SWAP_16");
 
   // Allow detection of fast FMA support.
   Builder.defineMacro("__FP_FAST_FMA", "1");
@@ -1061,7 +1062,7 @@ ParsedTargetAttr AArch64TargetInfo::parseTargetAttr(StringRef Features) const {
       else
         // Pushing the original feature string to give a sema error later on
         // when they get checked.
-        if (Feature.startswith("no"))
+        if (Feature.starts_with("no"))
           Features.push_back("-" + Feature.drop_front(2).str());
         else
           Features.push_back("+" + Feature.str());
@@ -1070,15 +1071,15 @@ ParsedTargetAttr AArch64TargetInfo::parseTargetAttr(StringRef Features) const {
 
   for (auto &Feature : AttrFeatures) {
     Feature = Feature.trim();
-    if (Feature.startswith("fpmath="))
+    if (Feature.starts_with("fpmath="))
       continue;
 
-    if (Feature.startswith("branch-protection=")) {
+    if (Feature.starts_with("branch-protection=")) {
       Ret.BranchProtection = Feature.split('=').second.trim();
       continue;
     }
 
-    if (Feature.startswith("arch=")) {
+    if (Feature.starts_with("arch=")) {
       if (FoundArch)
         Ret.Duplicate = "arch=";
       FoundArch = true;
@@ -1094,7 +1095,7 @@ ParsedTargetAttr AArch64TargetInfo::parseTargetAttr(StringRef Features) const {
       Ret.Features.push_back(AI->ArchFeature.str());
       // Add any extra features, after the +
       SplitAndAddFeatures(Split.second, Ret.Features);
-    } else if (Feature.startswith("cpu=")) {
+    } else if (Feature.starts_with("cpu=")) {
       if (!Ret.CPU.empty())
         Ret.Duplicate = "cpu=";
       else {
@@ -1105,14 +1106,14 @@ ParsedTargetAttr AArch64TargetInfo::parseTargetAttr(StringRef Features) const {
         Ret.CPU = Split.first;
         SplitAndAddFeatures(Split.second, Ret.Features);
       }
-    } else if (Feature.startswith("tune=")) {
+    } else if (Feature.starts_with("tune=")) {
       if (!Ret.Tune.empty())
         Ret.Duplicate = "tune=";
       else
         Ret.Tune = Feature.split("=").second.trim();
-    } else if (Feature.startswith("+")) {
+    } else if (Feature.starts_with("+")) {
       SplitAndAddFeatures(Feature, Ret.Features);
-    } else if (Feature.startswith("no-")) {
+    } else if (Feature.starts_with("no-")) {
       StringRef FeatureName =
           llvm::AArch64::getArchExtFeature(Feature.split("-").second);
       if (!FeatureName.empty())

clang/lib/Basic/Targets/AMDGPU.cpp

@@ -32,55 +32,56 @@ static const char *const DataLayoutStringR600 =
 
 static const char *const DataLayoutStringAMDGCN =
     "e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6:32:32"
-    "-p7:160:256:256:32-p8:128:128-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128"
+    "-p7:160:256:256:32-p8:128:128-p9:192:256:256:32-i64:64-v16:16-v24:32-v32:"
+    "32-v48:64-v96:128"
     "-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-S32-A5-G1"
-    "-ni:7:8";
+    "-ni:7:8:9";
 
 const LangASMap AMDGPUTargetInfo::AMDGPUDefIsGenMap = {
-    Generic,  // Default
-    Global,   // opencl_global
-    Local,    // opencl_local
-    Constant, // opencl_constant
-    Private,  // opencl_private
-    Generic,  // opencl_generic
-    Global,   // opencl_global_device
-    Global,   // opencl_global_host
-    Global,   // cuda_device
-    Constant, // cuda_constant
-    Local,    // cuda_shared
-    Global,   // sycl_global
-    Global,   // sycl_global_device
-    Global,   // sycl_global_host
-    Local,    // sycl_local
-    Private,  // sycl_private
-    Generic,  // ptr32_sptr
-    Generic,  // ptr32_uptr
-    Generic,  // ptr64
-    Generic,  // hlsl_groupshared
+    llvm::AMDGPUAS::FLAT_ADDRESS,     // Default
+    llvm::AMDGPUAS::GLOBAL_ADDRESS,   // opencl_global
+    llvm::AMDGPUAS::LOCAL_ADDRESS,    // opencl_local
+    llvm::AMDGPUAS::CONSTANT_ADDRESS, // opencl_constant
+    llvm::AMDGPUAS::PRIVATE_ADDRESS,  // opencl_private
+    llvm::AMDGPUAS::FLAT_ADDRESS,     // opencl_generic
+    llvm::AMDGPUAS::GLOBAL_ADDRESS,   // opencl_global_device
+    llvm::AMDGPUAS::GLOBAL_ADDRESS,   // opencl_global_host
+    llvm::AMDGPUAS::GLOBAL_ADDRESS,   // cuda_device
+    llvm::AMDGPUAS::CONSTANT_ADDRESS, // cuda_constant
+    llvm::AMDGPUAS::LOCAL_ADDRESS,    // cuda_shared
+    llvm::AMDGPUAS::GLOBAL_ADDRESS,   // sycl_global
+    llvm::AMDGPUAS::GLOBAL_ADDRESS,   // sycl_global_device
+    llvm::AMDGPUAS::GLOBAL_ADDRESS,   // sycl_global_host
+    llvm::AMDGPUAS::LOCAL_ADDRESS,    // sycl_local
+    llvm::AMDGPUAS::PRIVATE_ADDRESS,  // sycl_private
+    llvm::AMDGPUAS::FLAT_ADDRESS,     // ptr32_sptr
+    llvm::AMDGPUAS::FLAT_ADDRESS,     // ptr32_uptr
+    llvm::AMDGPUAS::FLAT_ADDRESS,     // ptr64
+    llvm::AMDGPUAS::FLAT_ADDRESS,     // hlsl_groupshared
 };
 
 const LangASMap AMDGPUTargetInfo::AMDGPUDefIsPrivMap = {
-    Private,  // Default
-    Global,   // opencl_global
-    Local,    // opencl_local
-    Constant, // opencl_constant
-    Private,  // opencl_private
-    Generic,  // opencl_generic
-    Global,   // opencl_global_device
-    Global,   // opencl_global_host
-    Global,   // cuda_device
-    Constant, // cuda_constant
-    Local,    // cuda_shared
+    llvm::AMDGPUAS::PRIVATE_ADDRESS,  // Default
+    llvm::AMDGPUAS::GLOBAL_ADDRESS,   // opencl_global
+    llvm::AMDGPUAS::LOCAL_ADDRESS,    // opencl_local
+    llvm::AMDGPUAS::CONSTANT_ADDRESS, // opencl_constant
+    llvm::AMDGPUAS::PRIVATE_ADDRESS,  // opencl_private
+    llvm::AMDGPUAS::FLAT_ADDRESS,     // opencl_generic
+    llvm::AMDGPUAS::GLOBAL_ADDRESS,   // opencl_global_device
+    llvm::AMDGPUAS::GLOBAL_ADDRESS,   // opencl_global_host
+    llvm::AMDGPUAS::GLOBAL_ADDRESS,   // cuda_device
+    llvm::AMDGPUAS::CONSTANT_ADDRESS, // cuda_constant
+    llvm::AMDGPUAS::LOCAL_ADDRESS,    // cuda_shared
     // SYCL address space values for this map are dummy
-    Generic, // sycl_global
-    Generic, // sycl_global_device
-    Generic, // sycl_global_host
-    Generic, // sycl_local
-    Generic, // sycl_private
-    Generic, // ptr32_sptr
-    Generic, // ptr32_uptr
-    Generic, // ptr64
-    Generic, // hlsl_groupshared
+    llvm::AMDGPUAS::FLAT_ADDRESS, // sycl_global
+    llvm::AMDGPUAS::FLAT_ADDRESS, // sycl_global_device
+    llvm::AMDGPUAS::FLAT_ADDRESS, // sycl_global_host
+    llvm::AMDGPUAS::FLAT_ADDRESS, // sycl_local
+    llvm::AMDGPUAS::FLAT_ADDRESS, // sycl_private
+    llvm::AMDGPUAS::FLAT_ADDRESS, // ptr32_sptr
+    llvm::AMDGPUAS::FLAT_ADDRESS, // ptr32_uptr
+    llvm::AMDGPUAS::FLAT_ADDRESS, // ptr64
+    llvm::AMDGPUAS::FLAT_ADDRESS, // hlsl_groupshared
 
 };
 } // namespace targets
@@ -279,7 +280,7 @@ void AMDGPUTargetInfo::getTargetDefines(const LangOptions &Opts,
     Builder.defineMacro(Twine("__") + Twine(CanonName) + Twine("__"));
     // Emit macros for gfx family e.g. gfx906 -> __GFX9__, gfx1030 -> __GFX10___
     if (isAMDGCN(getTriple())) {
-      assert(CanonName.startswith("gfx") && "Invalid amdgcn canonical name");
+      assert(CanonName.starts_with("gfx") && "Invalid amdgcn canonical name");
       Builder.defineMacro(Twine("__") + Twine(CanonName.drop_back(2).upper()) +
                           Twine("__"));
     }

clang/lib/Basic/Targets/AMDGPU.h

@@ -17,6 +17,7 @@
 #include "clang/Basic/TargetInfo.h"
 #include "clang/Basic/TargetOptions.h"
 #include "llvm/ADT/StringSet.h"
+#include "llvm/Support/AMDGPUAddrSpace.h"
 #include "llvm/Support/Compiler.h"
 #include "llvm/TargetParser/TargetParser.h"
 #include "llvm/TargetParser/Triple.h"
@@ -29,13 +30,6 @@ class LLVM_LIBRARY_VISIBILITY AMDGPUTargetInfo final : public TargetInfo {
 
   static const char *const GCCRegNames[];
 
-  enum AddrSpace {
-    Generic = 0,
-    Global = 1,
-    Local = 3,
-    Constant = 4,
-    Private = 5
-  };
   static const LangASMap AMDGPUDefIsGenMap;
   static const LangASMap AMDGPUDefIsPrivMap;
 
@@ -106,7 +100,8 @@ class LLVM_LIBRARY_VISIBILITY AMDGPUTargetInfo final : public TargetInfo {
       return 32;
     unsigned TargetAS = getTargetAddressSpace(AS);
 
-    if (TargetAS == Private || TargetAS == Local)
+    if (TargetAS == llvm::AMDGPUAS::PRIVATE_ADDRESS ||
+        TargetAS == llvm::AMDGPUAS::LOCAL_ADDRESS)
       return 32;
 
     return 64;
@@ -376,7 +371,7 @@ class LLVM_LIBRARY_VISIBILITY AMDGPUTargetInfo final : public TargetInfo {
   }
 
   std::optional<LangAS> getConstantAddressSpace() const override {
-    return getLangASFromTargetAS(Constant);
+    return getLangASFromTargetAS(llvm::AMDGPUAS::CONSTANT_ADDRESS);
   }
 
   const llvm::omp::GV &getGridValue() const override {
@@ -392,7 +387,7 @@ class LLVM_LIBRARY_VISIBILITY AMDGPUTargetInfo final : public TargetInfo {
 
   /// \returns Target specific vtbl ptr address space.
   unsigned getVtblPtrAddressSpace() const override {
-    return static_cast<unsigned>(Constant);
+    return static_cast<unsigned>(llvm::AMDGPUAS::CONSTANT_ADDRESS);
   }
 
   /// \returns If a target requires an address within a target specific address
@@ -405,9 +400,9 @@ class LLVM_LIBRARY_VISIBILITY AMDGPUTargetInfo final : public TargetInfo {
   getDWARFAddressSpace(unsigned AddressSpace) const override {
     const unsigned DWARF_Private = 1;
     const unsigned DWARF_Local = 2;
-    if (AddressSpace == Private) {
+    if (AddressSpace == llvm::AMDGPUAS::PRIVATE_ADDRESS) {
       return DWARF_Private;
-    } else if (AddressSpace == Local) {
+    } else if (AddressSpace == llvm::AMDGPUAS::LOCAL_ADDRESS) {
       return DWARF_Local;
     } else {
       return std::nullopt;

clang/lib/Basic/Targets/Mips.cpp

@@ -196,7 +196,7 @@ void MipsTargetInfo::getTargetDefines(const LangOptions &Opts,
   else
     Builder.defineMacro("_MIPS_ARCH_" + StringRef(CPU).upper());
 
-  if (StringRef(CPU).startswith("octeon"))
+  if (StringRef(CPU).starts_with("octeon"))
     Builder.defineMacro("__OCTEON__");
 
   if (CPU != "mips1") {

clang/lib/Basic/Targets/NVPTX.cpp

@@ -42,7 +42,7 @@ NVPTXTargetInfo::NVPTXTargetInfo(const llvm::Triple &Triple,
   PTXVersion = 32;
   for (const StringRef Feature : Opts.FeaturesAsWritten) {
     int PTXV;
-    if (!Feature.startswith("+ptx") ||
+    if (!Feature.starts_with("+ptx") ||
         Feature.drop_front(4).getAsInteger(10, PTXV))
       continue;
     PTXVersion = PTXV; // TODO: should it be max(PTXVersion, PTXV)?
@@ -262,11 +262,14 @@ void NVPTXTargetInfo::getTargetDefines(const LangOptions &Opts,
       case CudaArch::SM_89:
         return "890";
       case CudaArch::SM_90:
+      case CudaArch::SM_90a:
         return "900";
       }
       llvm_unreachable("unhandled CudaArch");
     }();
     Builder.defineMacro("__CUDA_ARCH__", CUDAArchCode);
+    if (GPU == CudaArch::SM_90a)
+      Builder.defineMacro("__CUDA_ARCH_FEAT_SM90_ALL", "1");
   }
 }
 

clang/lib/Basic/Targets/RISCV.cpp

@@ -131,7 +131,7 @@ static unsigned getVersionValue(unsigned MajorVersion, unsigned MinorVersion) {
 void RISCVTargetInfo::getTargetDefines(const LangOptions &Opts,
                                        MacroBuilder &Builder) const {
   Builder.defineMacro("__riscv");
-  bool Is64Bit = getTriple().getArch() == llvm::Triple::riscv64;
+  bool Is64Bit = getTriple().isRISCV64();
   Builder.defineMacro("__riscv_xlen", Is64Bit ? "64" : "32");
   StringRef CodeModel = getTargetOpts().CodeModel;
   unsigned FLen = ISAInfo->getFLen();
@@ -281,7 +281,7 @@ bool RISCVTargetInfo::initFeatureMap(
 
   unsigned XLen = 32;
 
-  if (getTriple().getArch() == llvm::Triple::riscv64) {
+  if (getTriple().isRISCV64()) {
     Features["64bit"] = true;
     XLen = 64;
   } else {
@@ -304,11 +304,18 @@ bool RISCVTargetInfo::initFeatureMap(
 
   // RISCVISAInfo makes implications for ISA features
   std::vector<std::string> ImpliedFeatures = (*ParseResult)->toFeatureVector();
-  // Add non-ISA features like `relax` and `save-restore` back
-  for (const std::string &Feature : NewFeaturesVec)
-    if (!llvm::is_contained(ImpliedFeatures, Feature))
-      ImpliedFeatures.push_back(Feature);
 
+  // parseFeatures normalizes the feature set by dropping any explicit
+  // negatives, and non-extension features.  We need to preserve the later
+  // for correctness and want to preserve the former for consistency.
+  for (auto &Feature : NewFeaturesVec) {
+     StringRef ExtName = Feature;
+     assert(ExtName.size() > 1 && (ExtName[0] == '+' || ExtName[0] == '-'));
+     ExtName = ExtName.drop_front(1); // Drop '+' or '-'
+     if (!llvm::is_contained(ImpliedFeatures, ("+" + ExtName).str()) &&
+         !llvm::is_contained(ImpliedFeatures, ("-" + ExtName).str()))
+       ImpliedFeatures.push_back(Feature);
+  }
   return TargetInfo::initFeatureMap(Features, Diags, CPU, ImpliedFeatures);
 }
 
@@ -336,7 +343,7 @@ RISCVTargetInfo::getVScaleRange(const LangOptions &LangOpts) const {
 
 /// Return true if has this feature, need to sync with handleTargetFeatures.
 bool RISCVTargetInfo::hasFeature(StringRef Feature) const {
-  bool Is64Bit = getTriple().getArch() == llvm::Triple::riscv64;
+  bool Is64Bit = getTriple().isRISCV64();
   auto Result = llvm::StringSwitch<std::optional<bool>>(Feature)
                     .Case("riscv", true)
                     .Case("riscv32", !Is64Bit)
@@ -347,10 +354,7 @@ bool RISCVTargetInfo::hasFeature(StringRef Feature) const {
   if (Result)
     return *Result;
 
-  if (ISAInfo->isSupportedExtensionFeature(Feature))
   return ISAInfo->hasExtension(Feature);
-
-  return false;
 }
 
 /// Perform initialization based on the user configured set of features.
@@ -430,14 +434,14 @@ ParsedTargetAttr RISCVTargetInfo::parseTargetAttr(StringRef Features) const {
     Feature = Feature.trim();
     StringRef AttrString = Feature.split("=").second.trim();
 
-    if (Feature.startswith("arch=")) {
+    if (Feature.starts_with("arch=")) {
       // Override last features
       Ret.Features.clear();
       if (FoundArch)
         Ret.Duplicate = "arch=";
       FoundArch = true;
 
-      if (AttrString.startswith("+")) {
+      if (AttrString.starts_with("+")) {
         // EXTENSION like arch=+v,+zbb
         SmallVector<StringRef, 1> Exts;
         AttrString.split(Exts, ",");
@@ -457,7 +461,7 @@ ParsedTargetAttr RISCVTargetInfo::parseTargetAttr(StringRef Features) const {
         // full-arch-string like arch=rv64gcv
         handleFullArchString(AttrString, Ret.Features);
       }
-    } else if (Feature.startswith("cpu=")) {
+    } else if (Feature.starts_with("cpu=")) {
       if (!Ret.CPU.empty())
         Ret.Duplicate = "cpu=";
 
@@ -471,7 +475,7 @@ ParsedTargetAttr RISCVTargetInfo::parseTargetAttr(StringRef Features) const {
           handleFullArchString(MarchFromCPU, Ret.Features);
         }
       }
-    } else if (Feature.startswith("tune=")) {
+    } else if (Feature.starts_with("tune=")) {
       if (!Ret.Tune.empty())
         Ret.Duplicate = "tune=";
 

clang/lib/Basic/Warnings.cpp

@@ -97,7 +97,7 @@ void clang::ProcessWarningOptions(DiagnosticsEngine &Diags,
       // Check to see if this warning starts with "no-", if so, this is a
       // negative form of the option.
       bool isPositive = true;
-      if (Opt.startswith("no-")) {
+      if (Opt.starts_with("no-")) {
         isPositive = false;
         Opt = Opt.substr(3);
       }
@@ -133,7 +133,7 @@ void clang::ProcessWarningOptions(DiagnosticsEngine &Diags,
       // table. It also has the "specifier" form of -Werror=foo. GCC supports
       // the deprecated -Werror-implicit-function-declaration which is used by
       // a few projects.
-      if (Opt.startswith("error")) {
+      if (Opt.starts_with("error")) {
         StringRef Specifier;
         if (Opt.size() > 5) {  // Specifier must be present.
           if (Opt[5] != '=' &&
@@ -162,7 +162,7 @@ void clang::ProcessWarningOptions(DiagnosticsEngine &Diags,
       }
 
       // -Wfatal-errors is yet another special case.
-      if (Opt.startswith("fatal-errors")) {
+      if (Opt.starts_with("fatal-errors")) {
         StringRef Specifier;
         if (Opt.size() != 12) {
           if ((Opt[12] != '=' && Opt[12] != '-') || Opt.size() == 13) {
@@ -204,7 +204,7 @@ void clang::ProcessWarningOptions(DiagnosticsEngine &Diags,
 
       // Check to see if this warning starts with "no-", if so, this is a
       // negative form of the option.
-      bool IsPositive = !Opt.startswith("no-");
+      bool IsPositive = !Opt.starts_with("no-");
       if (!IsPositive) Opt = Opt.substr(3);
 
       auto Severity = IsPositive ? diag::Severity::Remark

clang/lib/CodeGen/BackendUtil.cpp

@@ -45,6 +45,7 @@
 #include "llvm/Passes/PassBuilder.h"
 #include "llvm/Passes/PassPlugin.h"
 #include "llvm/Passes/StandardInstrumentations.h"
+#include "llvm/ProfileData/InstrProfCorrelator.h"
 #include "llvm/Support/BuryPointer.h"
 #include "llvm/Support/CommandLine.h"
 #include "llvm/Support/MemoryBuffer.h"
@@ -102,17 +103,21 @@ static cl::opt<bool> ClSanitizeOnOptimizerEarlyEP(
     "sanitizer-early-opt-ep", cl::Optional,
     cl::desc("Insert sanitizers on OptimizerEarlyEP."), cl::init(false));
 
+extern cl::opt<InstrProfCorrelator::ProfCorrelatorKind> ProfileCorrelate;
+
 // Re-link builtin bitcodes after optimization
 cl::opt<bool> ClRelinkBuiltinBitcodePostop(
     "relink-builtin-bitcode-postop", cl::Optional,
     cl::desc("Re-link builtin bitcodes after optimization."), cl::init(false));
-}
+} // namespace llvm
 
 namespace {
 
 // Default filename used for profile generation.
 std::string getDefaultProfileGenName() {
-  return DebugInfoCorrelate ? "default_%m.proflite" : "default_%m.profraw";
+  return DebugInfoCorrelate || ProfileCorrelate != InstrProfCorrelator::NONE
+             ? "default_%m.proflite"
+             : "default_%m.profraw";
 }
 
 class EmitAssemblyHelper {
@@ -204,7 +209,7 @@ class EmitAssemblyHelper {
   void EmitAssembly(BackendAction Action, std::unique_ptr<raw_pwrite_stream> OS,
                     BackendConsumer *BC);
 };
-}
+} // namespace
 
 static SanitizerCoverageOptions
 getSancovOptsFromCGOpts(const CodeGenOptions &CGOpts) {
@@ -982,7 +987,7 @@ void EmitAssemblyHelper::RunOptimizationPipeline(
             getInstrProfOptions(CodeGenOpts, LangOpts))
       PB.registerPipelineStartEPCallback(
           [Options](ModulePassManager &MPM, OptimizationLevel Level) {
-            MPM.addPass(InstrProfiling(*Options, false));
+            MPM.addPass(InstrProfilingLoweringPass(*Options, false));
           });
 
     // TODO: Consider passing the MemoryProfileOutput to the pass builder via

clang/lib/CodeGen/CGBuiltin.cpp

@@ -3410,9 +3410,16 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
                                    { Src0->getType(), Src1->getType() });
     return RValue::get(Builder.CreateCall(F, { Src0, Src1 }));
   }
+  case Builtin::BI__builtin_frexpl: {
+    // Linux PPC will not be adding additional PPCDoubleDouble support.
+    // WIP to switch default to IEEE long double. Will emit libcall for
+    // frexpl instead of legalizing this type in the BE.
+    if (&getTarget().getLongDoubleFormat() == &llvm::APFloat::PPCDoubleDouble())
+      break;
+    LLVM_FALLTHROUGH;
+  }
   case Builtin::BI__builtin_frexp:
   case Builtin::BI__builtin_frexpf:
-  case Builtin::BI__builtin_frexpl:
   case Builtin::BI__builtin_frexpf128:
   case Builtin::BI__builtin_frexpf16:
     return RValue::get(emitFrexpBuiltin(*this, E, Intrinsic::frexp));
@@ -9985,6 +9992,10 @@ CodeGenFunction::getSVEOverloadTypes(const SVETypeFlags &TypeFlags,
   if (TypeFlags.isOverloadCvt())
     return {Ops[0]->getType(), Ops.back()->getType()};
 
+  if (TypeFlags.isReductionQV() && !ResultType->isScalableTy() &&
+      ResultType->isVectorTy())
+    return {ResultType, Ops[1]->getType()};
+
   assert(TypeFlags.isOverloadDefault() && "Unexpected value for overloads");
   return {DefaultType};
 }

clang/lib/CodeGen/CGCUDANV.cpp

@@ -39,7 +39,7 @@ class CGNVCUDARuntime : public CGCUDARuntime {
 private:
   llvm::IntegerType *IntTy, *SizeTy;
   llvm::Type *VoidTy;
-  llvm::PointerType *CharPtrTy, *VoidPtrTy, *VoidPtrPtrTy;
+  llvm::PointerType *PtrTy;
 
   /// Convenience reference to LLVM Context
   llvm::LLVMContext &Context;
@@ -232,15 +232,12 @@ CGNVCUDARuntime::CGNVCUDARuntime(CodeGenModule &CGM)
   VoidTy = CGM.VoidTy;
   Zeros[0] = llvm::ConstantInt::get(SizeTy, 0);
   Zeros[1] = Zeros[0];
-
-  CharPtrTy = CGM.UnqualPtrTy;
-  VoidPtrTy = CGM.UnqualPtrTy;
-  VoidPtrPtrTy = CGM.UnqualPtrTy;
+  PtrTy = CGM.UnqualPtrTy;
 }
 
 llvm::FunctionCallee CGNVCUDARuntime::getSetupArgumentFn() const {
   // cudaError_t cudaSetupArgument(void *, size_t, size_t)
-  llvm::Type *Params[] = {VoidPtrTy, SizeTy, SizeTy};
+  llvm::Type *Params[] = {PtrTy, SizeTy, SizeTy};
   return CGM.CreateRuntimeFunction(
       llvm::FunctionType::get(IntTy, Params, false),
       addPrefixToName("SetupArgument"));
@@ -250,24 +247,24 @@ llvm::FunctionCallee CGNVCUDARuntime::getLaunchFn() const {
   if (CGM.getLangOpts().HIP) {
     // hipError_t hipLaunchByPtr(char *);
     return CGM.CreateRuntimeFunction(
-        llvm::FunctionType::get(IntTy, CharPtrTy, false), "hipLaunchByPtr");
+        llvm::FunctionType::get(IntTy, PtrTy, false), "hipLaunchByPtr");
   }
   // cudaError_t cudaLaunch(char *);
-  return CGM.CreateRuntimeFunction(
-      llvm::FunctionType::get(IntTy, CharPtrTy, false), "cudaLaunch");
+  return CGM.CreateRuntimeFunction(llvm::FunctionType::get(IntTy, PtrTy, false),
+                                   "cudaLaunch");
 }
 
 llvm::FunctionType *CGNVCUDARuntime::getRegisterGlobalsFnTy() const {
-  return llvm::FunctionType::get(VoidTy, VoidPtrPtrTy, false);
+  return llvm::FunctionType::get(VoidTy, PtrTy, false);
 }
 
 llvm::FunctionType *CGNVCUDARuntime::getCallbackFnTy() const {
-  return llvm::FunctionType::get(VoidTy, VoidPtrTy, false);
+  return llvm::FunctionType::get(VoidTy, PtrTy, false);
 }
 
 llvm::FunctionType *CGNVCUDARuntime::getRegisterLinkedBinaryFnTy() const {
-  llvm::Type *Params[] = {llvm::PointerType::getUnqual(Context), VoidPtrTy,
-                          VoidPtrTy, llvm::PointerType::getUnqual(Context)};
+  llvm::Type *Params[] = {llvm::PointerType::getUnqual(Context), PtrTy, PtrTy,
+                          llvm::PointerType::getUnqual(Context)};
   return llvm::FunctionType::get(VoidTy, Params, false);
 }
 
@@ -330,15 +327,15 @@ void CGNVCUDARuntime::emitDeviceStubBodyNew(CodeGenFunction &CGF,
   // args, allocate a single pointer so we still have a valid pointer to the
   // argument array that we can pass to runtime, even if it will be unused.
   Address KernelArgs = CGF.CreateTempAlloca(
-      VoidPtrTy, CharUnits::fromQuantity(16), "kernel_args",
+      PtrTy, CharUnits::fromQuantity(16), "kernel_args",
       llvm::ConstantInt::get(SizeTy, std::max<size_t>(1, Args.size())));
   // Store pointers to the arguments in a locally allocated launch_args.
   for (unsigned i = 0; i < Args.size(); ++i) {
     llvm::Value* VarPtr = CGF.GetAddrOfLocalVar(Args[i]).getPointer();
-    llvm::Value *VoidVarPtr = CGF.Builder.CreatePointerCast(VarPtr, VoidPtrTy);
+    llvm::Value *VoidVarPtr = CGF.Builder.CreatePointerCast(VarPtr, PtrTy);
     CGF.Builder.CreateDefaultAlignedStore(
         VoidVarPtr,
-        CGF.Builder.CreateConstGEP1_32(VoidPtrTy, KernelArgs.getPointer(), i));
+        CGF.Builder.CreateConstGEP1_32(PtrTy, KernelArgs.getPointer(), i));
   }
 
   llvm::BasicBlock *EndBlock = CGF.createBasicBlock("setup.end");
@@ -386,8 +383,7 @@ void CGNVCUDARuntime::emitDeviceStubBodyNew(CodeGenFunction &CGF,
       CGF.CreateMemTemp(Dim3Ty, CharUnits::fromQuantity(8), "block_dim");
   Address ShmemSize =
       CGF.CreateTempAlloca(SizeTy, CGM.getSizeAlign(), "shmem_size");
-  Address Stream =
-      CGF.CreateTempAlloca(VoidPtrTy, CGM.getPointerAlign(), "stream");
+  Address Stream = CGF.CreateTempAlloca(PtrTy, CGM.getPointerAlign(), "stream");
   llvm::FunctionCallee cudaPopConfigFn = CGM.CreateRuntimeFunction(
       llvm::FunctionType::get(IntTy,
                               {/*gridDim=*/GridDim.getType(),
@@ -402,8 +398,8 @@ void CGNVCUDARuntime::emitDeviceStubBodyNew(CodeGenFunction &CGF,
                                ShmemSize.getPointer(), Stream.getPointer()});
 
   // Emit the call to cudaLaunch
-  llvm::Value *Kernel = CGF.Builder.CreatePointerCast(
-      KernelHandles[CGF.CurFn->getName()], VoidPtrTy);
+  llvm::Value *Kernel =
+      CGF.Builder.CreatePointerCast(KernelHandles[CGF.CurFn->getName()], PtrTy);
   CallArgList LaunchKernelArgs;
   LaunchKernelArgs.add(RValue::get(Kernel),
                        cudaLaunchKernelFD->getParamDecl(0)->getType());
@@ -443,7 +439,7 @@ void CGNVCUDARuntime::emitDeviceStubBodyLegacy(CodeGenFunction &CGF,
     Offset = Offset.alignTo(TInfo.Align);
     llvm::Value *Args[] = {
         CGF.Builder.CreatePointerCast(CGF.GetAddrOfLocalVar(A).getPointer(),
-                                      VoidPtrTy),
+                                      PtrTy),
         llvm::ConstantInt::get(SizeTy, TInfo.Width.getQuantity()),
         llvm::ConstantInt::get(SizeTy, Offset.getQuantity()),
     };
@@ -458,8 +454,8 @@ void CGNVCUDARuntime::emitDeviceStubBodyLegacy(CodeGenFunction &CGF,
 
   // Emit the call to cudaLaunch
   llvm::FunctionCallee cudaLaunchFn = getLaunchFn();
-  llvm::Value *Arg = CGF.Builder.CreatePointerCast(
-      KernelHandles[CGF.CurFn->getName()], CharPtrTy);
+  llvm::Value *Arg =
+      CGF.Builder.CreatePointerCast(KernelHandles[CGF.CurFn->getName()], PtrTy);
   CGF.EmitRuntimeCallOrInvoke(cudaLaunchFn, Arg);
   CGF.EmitBranch(EndBlock);
 
@@ -537,11 +533,8 @@ llvm::Function *CGNVCUDARuntime::makeRegisterGlobalsFn() {
   // void __cudaRegisterFunction(void **, const char *, char *, const char *,
   //                             int, uint3*, uint3*, dim3*, dim3*, int*)
   llvm::Type *RegisterFuncParams[] = {
-      VoidPtrPtrTy, CharPtrTy,
-      CharPtrTy,    CharPtrTy,
-      IntTy,        VoidPtrTy,
-      VoidPtrTy,    VoidPtrTy,
-      VoidPtrTy,    llvm::PointerType::getUnqual(Context)};
+      PtrTy, PtrTy, PtrTy, PtrTy, IntTy,
+      PtrTy, PtrTy, PtrTy, PtrTy, llvm::PointerType::getUnqual(Context)};
   llvm::FunctionCallee RegisterFunc = CGM.CreateRuntimeFunction(
       llvm::FunctionType::get(IntTy, RegisterFuncParams, false),
       addUnderscoredPrefixToName("RegisterFunction"));
@@ -553,7 +546,7 @@ llvm::Function *CGNVCUDARuntime::makeRegisterGlobalsFn() {
   for (auto &&I : EmittedKernels) {
     llvm::Constant *KernelName =
         makeConstantString(getDeviceSideName(cast<NamedDecl>(I.D)));
-    llvm::Constant *NullPtr = llvm::ConstantPointerNull::get(VoidPtrTy);
+    llvm::Constant *NullPtr = llvm::ConstantPointerNull::get(PtrTy);
     llvm::Value *Args[] = {
         &GpuBinaryHandlePtr,
         KernelHandles[I.Kernel->getName()],
@@ -576,16 +569,15 @@ llvm::Function *CGNVCUDARuntime::makeRegisterGlobalsFn() {
 
   // void __cudaRegisterVar(void **, char *, char *, const char *,
   //                        int, int, int, int)
-  llvm::Type *RegisterVarParams[] = {VoidPtrPtrTy, CharPtrTy, CharPtrTy,
-                                     CharPtrTy,    IntTy,     VarSizeTy,
-                                     IntTy,        IntTy};
+  llvm::Type *RegisterVarParams[] = {PtrTy, PtrTy,     PtrTy, PtrTy,
+                                     IntTy, VarSizeTy, IntTy, IntTy};
   llvm::FunctionCallee RegisterVar = CGM.CreateRuntimeFunction(
       llvm::FunctionType::get(VoidTy, RegisterVarParams, false),
       addUnderscoredPrefixToName("RegisterVar"));
   // void __hipRegisterManagedVar(void **, char *, char *, const char *,
   //                              size_t, unsigned)
-  llvm::Type *RegisterManagedVarParams[] = {VoidPtrPtrTy, CharPtrTy, CharPtrTy,
-                                            CharPtrTy,    VarSizeTy, IntTy};
+  llvm::Type *RegisterManagedVarParams[] = {PtrTy, PtrTy,     PtrTy,
+                                            PtrTy, VarSizeTy, IntTy};
   llvm::FunctionCallee RegisterManagedVar = CGM.CreateRuntimeFunction(
       llvm::FunctionType::get(VoidTy, RegisterManagedVarParams, false),
       addUnderscoredPrefixToName("RegisterManagedVar"));
@@ -593,16 +585,13 @@ llvm::Function *CGNVCUDARuntime::makeRegisterGlobalsFn() {
   //                            const void **, const char *, int, int);
   llvm::FunctionCallee RegisterSurf = CGM.CreateRuntimeFunction(
       llvm::FunctionType::get(
-          VoidTy, {VoidPtrPtrTy, VoidPtrTy, CharPtrTy, CharPtrTy, IntTy, IntTy},
-          false),
+          VoidTy, {PtrTy, PtrTy, PtrTy, PtrTy, IntTy, IntTy}, false),
       addUnderscoredPrefixToName("RegisterSurface"));
   // void __cudaRegisterTexture(void **, const struct textureReference *,
   //                            const void **, const char *, int, int, int)
   llvm::FunctionCallee RegisterTex = CGM.CreateRuntimeFunction(
       llvm::FunctionType::get(
-          VoidTy,
-          {VoidPtrPtrTy, VoidPtrTy, CharPtrTy, CharPtrTy, IntTy, IntTy, IntTy},
-          false),
+          VoidTy, {PtrTy, PtrTy, PtrTy, PtrTy, IntTy, IntTy, IntTy}, false),
       addUnderscoredPrefixToName("RegisterTexture"));
   for (auto &&Info : DeviceVars) {
     llvm::GlobalVariable *Var = Info.Var;
@@ -713,11 +702,11 @@ llvm::Function *CGNVCUDARuntime::makeModuleCtorFunction() {
 
   // void ** __{cuda|hip}RegisterFatBinary(void *);
   llvm::FunctionCallee RegisterFatbinFunc = CGM.CreateRuntimeFunction(
-      llvm::FunctionType::get(VoidPtrPtrTy, VoidPtrTy, false),
+      llvm::FunctionType::get(PtrTy, PtrTy, false),
       addUnderscoredPrefixToName("RegisterFatBinary"));
   // struct { int magic, int version, void * gpu_binary, void * dont_care };
   llvm::StructType *FatbinWrapperTy =
-      llvm::StructType::get(IntTy, IntTy, VoidPtrTy, VoidPtrTy);
+      llvm::StructType::get(IntTy, IntTy, PtrTy, PtrTy);
 
   // Register GPU binary with the CUDA runtime, store returned handle in a
   // global variable and save a reference in GpuBinaryHandle to be cleaned up
@@ -813,7 +802,7 @@ llvm::Function *CGNVCUDARuntime::makeModuleCtorFunction() {
   // Data.
   Values.add(FatBinStr);
   // Unused in fatbin v1.
-  Values.add(llvm::ConstantPointerNull::get(VoidPtrTy));
+  Values.add(llvm::ConstantPointerNull::get(PtrTy));
   llvm::GlobalVariable *FatbinWrapper = Values.finishAndCreateGlobal(
       addUnderscoredPrefixToName("_fatbin_wrapper"), CGM.getPointerAlign(),
       /*constant*/ true);
@@ -836,9 +825,8 @@ llvm::Function *CGNVCUDARuntime::makeModuleCtorFunction() {
     // The name, size, and initialization pattern of this variable is part
     // of HIP ABI.
     GpuBinaryHandle = new llvm::GlobalVariable(
-        TheModule, VoidPtrPtrTy, /*isConstant=*/false,
-        Linkage,
-        /*Initializer=*/llvm::ConstantPointerNull::get(VoidPtrPtrTy),
+        TheModule, PtrTy, /*isConstant=*/false, Linkage,
+        /*Initializer=*/llvm::ConstantPointerNull::get(PtrTy),
         "__hip_gpubin_handle");
     if (Linkage == llvm::GlobalValue::LinkOnceAnyLinkage)
       GpuBinaryHandle->setComdat(
@@ -848,7 +836,7 @@ llvm::Function *CGNVCUDARuntime::makeModuleCtorFunction() {
     if (Linkage != llvm::GlobalValue::InternalLinkage)
       GpuBinaryHandle->setVisibility(llvm::GlobalValue::HiddenVisibility);
     Address GpuBinaryAddr(
-        GpuBinaryHandle, VoidPtrPtrTy,
+        GpuBinaryHandle, PtrTy,
         CharUnits::fromQuantity(GpuBinaryHandle->getAlignment()));
     {
       auto *HandleValue = CtorBuilder.CreateLoad(GpuBinaryAddr);
@@ -880,8 +868,8 @@ llvm::Function *CGNVCUDARuntime::makeModuleCtorFunction() {
     llvm::CallInst *RegisterFatbinCall =
         CtorBuilder.CreateCall(RegisterFatbinFunc, FatbinWrapper);
     GpuBinaryHandle = new llvm::GlobalVariable(
-        TheModule, VoidPtrPtrTy, false, llvm::GlobalValue::InternalLinkage,
-        llvm::ConstantPointerNull::get(VoidPtrPtrTy), "__cuda_gpubin_handle");
+        TheModule, PtrTy, false, llvm::GlobalValue::InternalLinkage,
+        llvm::ConstantPointerNull::get(PtrTy), "__cuda_gpubin_handle");
     GpuBinaryHandle->setAlignment(CGM.getPointerAlign().getAsAlign());
     CtorBuilder.CreateAlignedStore(RegisterFatbinCall, GpuBinaryHandle,
                                    CGM.getPointerAlign());
@@ -895,7 +883,7 @@ llvm::Function *CGNVCUDARuntime::makeModuleCtorFunction() {
                            CudaFeature::CUDA_USES_FATBIN_REGISTER_END)) {
       // void __cudaRegisterFatBinaryEnd(void **);
       llvm::FunctionCallee RegisterFatbinEndFunc = CGM.CreateRuntimeFunction(
-          llvm::FunctionType::get(VoidTy, VoidPtrPtrTy, false),
+          llvm::FunctionType::get(VoidTy, PtrTy, false),
           "__cudaRegisterFatBinaryEnd");
       CtorBuilder.CreateCall(RegisterFatbinEndFunc, RegisterFatbinCall);
     }
@@ -967,7 +955,7 @@ llvm::Function *CGNVCUDARuntime::makeModuleDtorFunction() {
 
   // void __cudaUnregisterFatBinary(void ** handle);
   llvm::FunctionCallee UnregisterFatbinFunc = CGM.CreateRuntimeFunction(
-      llvm::FunctionType::get(VoidTy, VoidPtrPtrTy, false),
+      llvm::FunctionType::get(VoidTy, PtrTy, false),
       addUnderscoredPrefixToName("UnregisterFatBinary"));
 
   llvm::Function *ModuleDtorFunc = llvm::Function::Create(

clang/lib/CodeGen/CGCall.cpp

@@ -5609,7 +5609,7 @@ RValue CodeGenFunction::EmitCall(const CGFunctionInfo &CallInfo,
     EmitBlock(Cont);
   }
   if (CI->getCalledFunction() && CI->getCalledFunction()->hasName() &&
-      CI->getCalledFunction()->getName().startswith("_Z4sqrt")) {
+      CI->getCalledFunction()->getName().starts_with("_Z4sqrt")) {
     SetSqrtFPAccuracy(CI);
   }
   if (callOrInvoke)

clang/lib/CodeGen/CGCall.h

@@ -20,6 +20,7 @@
 #include "clang/AST/CanonicalType.h"
 #include "clang/AST/GlobalDecl.h"
 #include "clang/AST/Type.h"
+#include "llvm/ADT/STLForwardCompat.h"
 #include "llvm/IR/Value.h"
 
 namespace llvm {
@@ -406,15 +407,13 @@ enum class FnInfoOpts {
 };
 
 inline FnInfoOpts operator|(FnInfoOpts A, FnInfoOpts B) {
-  return static_cast<FnInfoOpts>(
-      static_cast<std::underlying_type_t<FnInfoOpts>>(A) |
-      static_cast<std::underlying_type_t<FnInfoOpts>>(B));
+  return static_cast<FnInfoOpts>(llvm::to_underlying(A) |
+                                 llvm::to_underlying(B));
 }
 
 inline FnInfoOpts operator&(FnInfoOpts A, FnInfoOpts B) {
-  return static_cast<FnInfoOpts>(
-      static_cast<std::underlying_type_t<FnInfoOpts>>(A) &
-      static_cast<std::underlying_type_t<FnInfoOpts>>(B));
+  return static_cast<FnInfoOpts>(llvm::to_underlying(A) &
+                                 llvm::to_underlying(B));
 }
 
 inline FnInfoOpts operator|=(FnInfoOpts A, FnInfoOpts B) {

clang/lib/CodeGen/CGDebugInfo.cpp

@@ -554,15 +554,17 @@ void CGDebugInfo::CreateCompileUnit() {
     // If the main file name provided is identical to the input file name, and
     // if the input file is a preprocessed source, use the module name for
     // debug info. The module name comes from the name specified in the first
-    // linemarker if the input is a preprocessed source.
+    // linemarker if the input is a preprocessed source. In this case we don't
+    // know the content to compute a checksum.
     if (MainFile->getName() == MainFileName &&
         FrontendOptions::getInputKindForExtension(
             MainFile->getName().rsplit('.').second)
-            .isPreprocessed())
+            .isPreprocessed()) {
       MainFileName = CGM.getModule().getName().str();
-
+    } else {
       CSKind = computeChecksum(SM.getMainFileID(), Checksum);
     }
+  }
 
   llvm::dwarf::SourceLanguage LangTag;
   if (LO.CPlusPlus) {
@@ -635,7 +637,8 @@ void CGDebugInfo::CreateCompileUnit() {
     Sysroot = CGM.getHeaderSearchOpts().Sysroot;
     auto B = llvm::sys::path::rbegin(Sysroot);
     auto E = llvm::sys::path::rend(Sysroot);
-    auto It = std::find_if(B, E, [](auto SDK) { return SDK.endswith(".sdk"); });
+    auto It =
+        std::find_if(B, E, [](auto SDK) { return SDK.ends_with(".sdk"); });
     if (It != E)
       SDK = *It;
   }
@@ -2883,7 +2886,7 @@ llvm::DIModule *CGDebugInfo::getOrCreateModuleRef(ASTSourceDescriptor Mod,
   // clang::Module object, but it won't actually be built or imported; it will
   // be textual.
   if (CreateSkeletonCU && IsRootModule && Mod.getASTFile().empty() && M)
-    assert(StringRef(M->Name).startswith(CGM.getLangOpts().ModuleName) &&
+    assert(StringRef(M->Name).starts_with(CGM.getLangOpts().ModuleName) &&
            "clang module without ASTFile must be specified by -fmodule-name");
 
   // Return a StringRef to the remapped Path.
@@ -4247,7 +4250,7 @@ void CGDebugInfo::emitFunctionStart(GlobalDecl GD, SourceLocation Loc,
 
     Flags |= llvm::DINode::FlagPrototyped;
   }
-  if (Name.startswith("\01"))
+  if (Name.starts_with("\01"))
     Name = Name.substr(1);
 
   assert((!D || !isa<VarDecl>(D) ||

clang/lib/CodeGen/CGException.cpp

@@ -277,7 +277,7 @@ static bool LandingPadHasOnlyCXXUses(llvm::LandingPadInst *LPI) {
       if (llvm::GlobalVariable *GV = dyn_cast<llvm::GlobalVariable>(Val))
         // ObjC EH selector entries are always global variables with
         // names starting like this.
-        if (GV->getName().startswith("OBJC_EHTYPE"))
+        if (GV->getName().starts_with("OBJC_EHTYPE"))
           return false;
     } else {
       // Check if any of the filter values have the ObjC prefix.
@@ -288,7 +288,7 @@ static bool LandingPadHasOnlyCXXUses(llvm::LandingPadInst *LPI) {
             cast<llvm::GlobalVariable>((*II)->stripPointerCasts()))
           // ObjC EH selector entries are always global variables with
           // names starting like this.
-          if (GV->getName().startswith("OBJC_EHTYPE"))
+          if (GV->getName().starts_with("OBJC_EHTYPE"))
             return false;
       }
     }
@@ -1917,7 +1917,7 @@ void CodeGenFunction::EmitCapturedLocals(CodeGenFunction &ParentCGF,
         const VarDecl *D = cast<VarDecl>(I.first);
         if (isa<ImplicitParamDecl>(D) &&
             D->getType() == getContext().VoidPtrTy) {
-          assert(D->getName().startswith("frame_pointer"));
+          assert(D->getName().starts_with("frame_pointer"));
           FramePtrAddrAlloca = cast<llvm::AllocaInst>(I.second.getPointer());
           break;
         }

clang/lib/CodeGen/CGExpr.cpp

@@ -431,7 +431,7 @@ static Address createReferenceTemporary(CodeGenFunction &CGF,
 
 /// Helper method to check if the underlying ABI is AAPCS
 static bool isAAPCS(const TargetInfo &TargetInfo) {
-  return TargetInfo.getABI().startswith("aapcs");
+  return TargetInfo.getABI().starts_with("aapcs");
 }
 
 LValue CodeGenFunction::
@@ -3156,7 +3156,7 @@ LValue CodeGenFunction::EmitPredefinedLValue(const PredefinedExpr *E) {
   auto SL = E->getFunctionName();
   assert(SL != nullptr && "No StringLiteral name in PredefinedExpr");
   StringRef FnName = CurFn->getName();
-  if (FnName.startswith("\01"))
+  if (FnName.starts_with("\01"))
     FnName = FnName.substr(1);
   StringRef NameItems[] = {
       PredefinedExpr::getIdentKindName(E->getIdentKind()), FnName};

clang/lib/CodeGen/CGExprComplex.cpp

@@ -275,6 +275,10 @@ class ComplexExprEmitter
   ComplexPairTy EmitBinSub(const BinOpInfo &Op);
   ComplexPairTy EmitBinMul(const BinOpInfo &Op);
   ComplexPairTy EmitBinDiv(const BinOpInfo &Op);
+  ComplexPairTy EmitAlgebraicDiv(llvm::Value *A, llvm::Value *B, llvm::Value *C,
+                                 llvm::Value *D);
+  ComplexPairTy EmitRangeReductionDiv(llvm::Value *A, llvm::Value *B,
+                                      llvm::Value *C, llvm::Value *D);
 
   ComplexPairTy EmitComplexBinOpLibCall(StringRef LibCallName,
                                         const BinOpInfo &Op);
@@ -781,6 +785,10 @@ ComplexPairTy ComplexExprEmitter::EmitBinMul(const BinOpInfo &Op) {
       ResR = Builder.CreateFSub(AC, BD, "mul_r");
       ResI = Builder.CreateFAdd(AD, BC, "mul_i");
 
+      if (Op.FPFeatures.getComplexRange() == LangOptions::CX_Limited ||
+          Op.FPFeatures.getComplexRange() == LangOptions::CX_Fortran)
+        return ComplexPairTy(ResR, ResI);
+
       // Emit the test for the real part becoming NaN and create a branch to
       // handle it. We test for NaN by comparing the number to itself.
       Value *IsRNaN = Builder.CreateFCmpUNO(ResR, ResR, "isnan_cmp");
@@ -846,14 +854,132 @@ ComplexPairTy ComplexExprEmitter::EmitBinMul(const BinOpInfo &Op) {
   return ComplexPairTy(ResR, ResI);
 }
 
+ComplexPairTy ComplexExprEmitter::EmitAlgebraicDiv(llvm::Value *LHSr,
+                                                   llvm::Value *LHSi,
+                                                   llvm::Value *RHSr,
+                                                   llvm::Value *RHSi) {
+  // (a+ib) / (c+id) = ((ac+bd)/(cc+dd)) + i((bc-ad)/(cc+dd))
+  llvm::Value *DSTr, *DSTi;
+
+  llvm::Value *AC = Builder.CreateFMul(LHSr, RHSr); // a*c
+  llvm::Value *BD = Builder.CreateFMul(LHSi, RHSi); // b*d
+  llvm::Value *ACpBD = Builder.CreateFAdd(AC, BD);  // ac+bd
+
+  llvm::Value *CC = Builder.CreateFMul(RHSr, RHSr); // c*c
+  llvm::Value *DD = Builder.CreateFMul(RHSi, RHSi); // d*d
+  llvm::Value *CCpDD = Builder.CreateFAdd(CC, DD);  // cc+dd
+
+  llvm::Value *BC = Builder.CreateFMul(LHSi, RHSr); // b*c
+  llvm::Value *AD = Builder.CreateFMul(LHSr, RHSi); // a*d
+  llvm::Value *BCmAD = Builder.CreateFSub(BC, AD);  // bc-ad
+
+  DSTr = Builder.CreateFDiv(ACpBD, CCpDD);
+  DSTi = Builder.CreateFDiv(BCmAD, CCpDD);
+  return ComplexPairTy(DSTr, DSTi);
+}
+
+// EmitFAbs - Emit a call to @llvm.fabs.
+static llvm::Value *EmitllvmFAbs(CodeGenFunction &CGF, llvm::Value *Value) {
+  llvm::Function *Func =
+      CGF.CGM.getIntrinsic(llvm::Intrinsic::fabs, Value->getType());
+  llvm::Value *Call = CGF.Builder.CreateCall(Func, Value);
+  return Call;
+}
+
+// EmitRangeReductionDiv - Implements Smith's algorithm for complex division.
+// SMITH, R. L. Algorithm 116: Complex division. Commun. ACM 5, 8 (1962).
+ComplexPairTy ComplexExprEmitter::EmitRangeReductionDiv(llvm::Value *LHSr,
+                                                        llvm::Value *LHSi,
+                                                        llvm::Value *RHSr,
+                                                        llvm::Value *RHSi) {
+  // (a + ib) / (c + id) = (e + if)
+  llvm::Value *FAbsRHSr = EmitllvmFAbs(CGF, RHSr); // |c|
+  llvm::Value *FAbsRHSi = EmitllvmFAbs(CGF, RHSi); // |d|
+  // |c| >= |d|
+  llvm::Value *IsR = Builder.CreateFCmpUGT(FAbsRHSr, FAbsRHSi, "abs_cmp");
+
+  llvm::BasicBlock *TrueBB =
+      CGF.createBasicBlock("abs_rhsr_greater_or_equal_abs_rhsi");
+  llvm::BasicBlock *FalseBB =
+      CGF.createBasicBlock("abs_rhsr_less_than_abs_rhsi");
+  llvm::BasicBlock *ContBB = CGF.createBasicBlock("complex_div");
+  Builder.CreateCondBr(IsR, TrueBB, FalseBB);
+
+  CGF.EmitBlock(TrueBB);
+  // abs(c) >= abs(d)
+  // r = d/c
+  // tmp = c + rd
+  // e = (a + br)/tmp
+  // f = (b - ar)/tmp
+  llvm::Value *DdC = Builder.CreateFDiv(RHSi, RHSr); // r=d/c
+
+  llvm::Value *RD = Builder.CreateFMul(DdC, RHSi);  // rd
+  llvm::Value *CpRD = Builder.CreateFAdd(RHSr, RD); // tmp=c+rd
+
+  llvm::Value *T3 = Builder.CreateFMul(LHSi, DdC);   // br
+  llvm::Value *T4 = Builder.CreateFAdd(LHSr, T3);    // a+br
+  llvm::Value *DSTTr = Builder.CreateFDiv(T4, CpRD); // (a+br)/tmp
+
+  llvm::Value *T5 = Builder.CreateFMul(LHSr, DdC);   // ar
+  llvm::Value *T6 = Builder.CreateFSub(LHSi, T5);    // b-ar
+  llvm::Value *DSTTi = Builder.CreateFDiv(T6, CpRD); // (b-ar)/tmp
+  Builder.CreateBr(ContBB);
+
+  CGF.EmitBlock(FalseBB);
+  // abs(c) < abs(d)
+  // r = c/d
+  // tmp = d + rc
+  // e = (ar + b)/tmp
+  // f = (br - a)/tmp
+  llvm::Value *CdD = Builder.CreateFDiv(RHSr, RHSi); // r=c/d
+
+  llvm::Value *RC = Builder.CreateFMul(CdD, RHSr);  // rc
+  llvm::Value *DpRC = Builder.CreateFAdd(RHSi, RC); // tmp=d+rc
+
+  llvm::Value *T7 = Builder.CreateFMul(LHSr, RC);    // ar
+  llvm::Value *T8 = Builder.CreateFAdd(T7, LHSi);    // ar+b
+  llvm::Value *DSTFr = Builder.CreateFDiv(T8, DpRC); // (ar+b)/tmp
+
+  llvm::Value *T9 = Builder.CreateFMul(LHSi, CdD);    // br
+  llvm::Value *T10 = Builder.CreateFSub(T9, LHSr);    // br-a
+  llvm::Value *DSTFi = Builder.CreateFDiv(T10, DpRC); // (br-a)/tmp
+  Builder.CreateBr(ContBB);
+
+  // Phi together the computation paths.
+  CGF.EmitBlock(ContBB);
+  llvm::PHINode *VALr = Builder.CreatePHI(DSTTr->getType(), 2);
+  VALr->addIncoming(DSTTr, TrueBB);
+  VALr->addIncoming(DSTFr, FalseBB);
+  llvm::PHINode *VALi = Builder.CreatePHI(DSTTi->getType(), 2);
+  VALi->addIncoming(DSTTi, TrueBB);
+  VALi->addIncoming(DSTFi, FalseBB);
+  return ComplexPairTy(VALr, VALi);
+}
+
 // See C11 Annex G.5.1 for the semantics of multiplicative operators on complex
 // typed values.
 ComplexPairTy ComplexExprEmitter::EmitBinDiv(const BinOpInfo &Op) {
   llvm::Value *LHSr = Op.LHS.first, *LHSi = Op.LHS.second;
   llvm::Value *RHSr = Op.RHS.first, *RHSi = Op.RHS.second;
-
   llvm::Value *DSTr, *DSTi;
   if (LHSr->getType()->isFloatingPointTy()) {
+    CodeGenFunction::CGFPOptionsRAII FPOptsRAII(CGF, Op.FPFeatures);
+    if (!RHSi) {
+      assert(LHSi && "Can have at most one non-complex operand!");
+
+      DSTr = Builder.CreateFDiv(LHSr, RHSr);
+      DSTi = Builder.CreateFDiv(LHSi, RHSr);
+      return ComplexPairTy(DSTr, DSTi);
+    }
+    llvm::Value *OrigLHSi = LHSi;
+    if (!LHSi)
+      LHSi = llvm::Constant::getNullValue(RHSi->getType());
+    if (Op.FPFeatures.getComplexRange() == LangOptions::CX_Fortran)
+      return EmitRangeReductionDiv(LHSr, LHSi, RHSr, RHSi);
+    else if (Op.FPFeatures.getComplexRange() == LangOptions::CX_Limited)
+      return EmitAlgebraicDiv(LHSr, LHSi, RHSr, RHSi);
+    else if (!CGF.getLangOpts().FastMath) {
+      LHSi = OrigLHSi;
       // If we have a complex operand on the RHS and FastMath is not allowed, we
       // delegate to a libcall to handle all of the complexities and minimize
       // underflow/overflow cases. When FastMath is allowed we construct the
@@ -861,8 +987,6 @@ ComplexPairTy ComplexExprEmitter::EmitBinDiv(const BinOpInfo &Op) {
       //
       // FIXME: We would be able to avoid the libcall in many places if we
       // supported imaginary types in addition to complex types.
-    CodeGenFunction::CGFPOptionsRAII FPOptsRAII(CGF, Op.FPFeatures);
-    if (RHSi && !CGF.getLangOpts().FastMath) {
       BinOpInfo LibCallOp = Op;
       // If LHS was a real, supply a null imaginary part.
       if (!LHSi)
@@ -884,30 +1008,8 @@ ComplexPairTy ComplexExprEmitter::EmitBinDiv(const BinOpInfo &Op) {
       case llvm::Type::FP128TyID:
         return EmitComplexBinOpLibCall("__divtc3", LibCallOp);
       }
-    } else if (RHSi) {
-      if (!LHSi)
-        LHSi = llvm::Constant::getNullValue(RHSi->getType());
-
-      // (a+ib) / (c+id) = ((ac+bd)/(cc+dd)) + i((bc-ad)/(cc+dd))
-      llvm::Value *AC = Builder.CreateFMul(LHSr, RHSr); // a*c
-      llvm::Value *BD = Builder.CreateFMul(LHSi, RHSi); // b*d
-      llvm::Value *ACpBD = Builder.CreateFAdd(AC, BD); // ac+bd
-
-      llvm::Value *CC = Builder.CreateFMul(RHSr, RHSr); // c*c
-      llvm::Value *DD = Builder.CreateFMul(RHSi, RHSi); // d*d
-      llvm::Value *CCpDD = Builder.CreateFAdd(CC, DD); // cc+dd
-
-      llvm::Value *BC = Builder.CreateFMul(LHSi, RHSr); // b*c
-      llvm::Value *AD = Builder.CreateFMul(LHSr, RHSi); // a*d
-      llvm::Value *BCmAD = Builder.CreateFSub(BC, AD); // bc-ad
-
-      DSTr = Builder.CreateFDiv(ACpBD, CCpDD);
-      DSTi = Builder.CreateFDiv(BCmAD, CCpDD);
     } else {
-      assert(LHSi && "Can have at most one non-complex operand!");
-
-      DSTr = Builder.CreateFDiv(LHSr, RHSr);
-      DSTi = Builder.CreateFDiv(LHSi, RHSr);
+      return EmitAlgebraicDiv(LHSr, LHSi, RHSr, RHSi);
     }
   } else {
     assert(Op.LHS.second && Op.RHS.second &&

clang/lib/CodeGen/CGHLSLRuntime.cpp

@@ -184,7 +184,8 @@ void CGHLSLRuntime::finishCodeGen() {
                                       : llvm::hlsl::ResourceKind::TBuffer;
     std::string TyName =
         Buf.Name.str() + (Buf.IsCBuffer ? ".cb." : ".tb.") + "ty";
-    addBufferResourceAnnotation(GV, TyName, RC, RK, Buf.Binding);
+    addBufferResourceAnnotation(GV, TyName, RC, RK, /*IsROV=*/false,
+                                Buf.Binding);
   }
 }
 
@@ -196,6 +197,7 @@ void CGHLSLRuntime::addBufferResourceAnnotation(llvm::GlobalVariable *GV,
                                                 llvm::StringRef TyName,
                                                 llvm::hlsl::ResourceClass RC,
                                                 llvm::hlsl::ResourceKind RK,
+                                                bool IsROV,
                                                 BufferResBinding &Binding) {
   llvm::Module &M = CGM.getModule();
 
@@ -219,7 +221,7 @@ void CGHLSLRuntime::addBufferResourceAnnotation(llvm::GlobalVariable *GV,
          "ResourceMD must have been set by the switch above.");
 
   llvm::hlsl::FrontendResource Res(
-      GV, TyName, RK, Binding.Reg.value_or(UINT_MAX), Binding.Space);
+      GV, TyName, RK, IsROV, Binding.Reg.value_or(UINT_MAX), Binding.Space);
   ResourceMD->addOperand(Res.getMetadata());
 }
 
@@ -236,10 +238,11 @@ void CGHLSLRuntime::annotateHLSLResource(const VarDecl *D, GlobalVariable *GV) {
 
   llvm::hlsl::ResourceClass RC = Attr->getResourceClass();
   llvm::hlsl::ResourceKind RK = Attr->getResourceKind();
+  bool IsROV = Attr->getIsROV();
 
   QualType QT(Ty, 0);
   BufferResBinding Binding(D->getAttr<HLSLResourceBindingAttr>());
-  addBufferResourceAnnotation(GV, QT.getAsString(), RC, RK, Binding);
+  addBufferResourceAnnotation(GV, QT.getAsString(), RC, RK, IsROV, Binding);
 }
 
 CGHLSLRuntime::BufferResBinding::BufferResBinding(

clang/lib/CodeGen/CGHLSLRuntime.h

@@ -92,7 +92,7 @@ class CGHLSLRuntime {
   void addBufferResourceAnnotation(llvm::GlobalVariable *GV,
                                    llvm::StringRef TyName,
                                    llvm::hlsl::ResourceClass RC,
-                                   llvm::hlsl::ResourceKind RK,
+                                   llvm::hlsl::ResourceKind RK, bool IsROV,
                                    BufferResBinding &Binding);
   void addConstant(VarDecl *D, Buffer &CB);
   void addBufferDecls(const DeclContext *DC, Buffer &CB);

clang/lib/CodeGen/CGObjCMac.cpp

@@ -1850,7 +1850,7 @@ static bool hasObjCExceptionAttribute(ASTContext &Context,
 static llvm::GlobalValue::LinkageTypes
 getLinkageTypeForObjCMetadata(CodeGenModule &CGM, StringRef Section) {
   if (CGM.getTriple().isOSBinFormatMachO() &&
-      (Section.empty() || Section.startswith("__DATA")))
+      (Section.empty() || Section.starts_with("__DATA")))
     return llvm::GlobalValue::InternalLinkage;
   return llvm::GlobalValue::PrivateLinkage;
 }
@@ -6162,7 +6162,7 @@ void CGObjCNonFragileABIMac::AddModuleClassList(
   // Section name is obtained by calling GetSectionName, which returns
   // sections in the __DATA segment on MachO.
   assert((!CGM.getTriple().isOSBinFormatMachO() ||
-          SectionName.startswith("__DATA")) &&
+          SectionName.starts_with("__DATA")) &&
          "SectionName expected to start with __DATA on MachO");
   llvm::GlobalVariable *GV = new llvm::GlobalVariable(
       CGM.getModule(), Init->getType(), false,

clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp

@@ -3483,6 +3483,7 @@ void CGOpenMPRuntimeGPU::processRequiresDirective(
       case CudaArch::SM_87:
       case CudaArch::SM_89:
       case CudaArch::SM_90:
+      case CudaArch::SM_90a:
       case CudaArch::GFX600:
       case CudaArch::GFX601:
       case CudaArch::GFX602:

clang/lib/CodeGen/CGRecordLayoutBuilder.cpp

@@ -111,7 +111,7 @@ struct CGRecordLowering {
 
   /// Helper function to check if we are targeting AAPCS.
   bool isAAPCS() const {
-    return Context.getTargetInfo().getABI().startswith("aapcs");
+    return Context.getTargetInfo().getABI().starts_with("aapcs");
   }
 
   /// Helper function to check if the target machine is BigEndian.

clang/lib/CodeGen/CGStmt.cpp

@@ -2548,7 +2548,7 @@ void CodeGenFunction::EmitAsmStmt(const AsmStmt &S) {
       ResultRegQualTys.push_back(QTy);
       ResultRegDests.push_back(Dest);
 
-      bool IsFlagReg = llvm::StringRef(OutputConstraint).startswith("{@cc");
+      bool IsFlagReg = llvm::StringRef(OutputConstraint).starts_with("{@cc");
       ResultRegIsFlagReg.push_back(IsFlagReg);
 
       llvm::Type *Ty = ConvertTypeForMem(QTy);

clang/lib/CodeGen/CodeGenAction.cpp

@@ -1139,7 +1139,7 @@ CodeGenAction::loadModule(MemoryBufferRef MBRef) {
 
   // Strip off a leading diagnostic code if there is one.
   StringRef Msg = Err.getMessage();
-  if (Msg.startswith("error: "))
+  if (Msg.starts_with("error: "))
     Msg = Msg.substr(7);
 
   unsigned DiagID =

clang/lib/CodeGen/CodeGenModule.cpp

@@ -225,9 +225,9 @@ createTargetCodeGenInfo(CodeGenModule &CGM) {
     StringRef ABIStr = Target.getABI();
     unsigned XLen = Target.getPointerWidth(LangAS::Default);
     unsigned ABIFLen = 0;
-    if (ABIStr.endswith("f"))
+    if (ABIStr.ends_with("f"))
       ABIFLen = 32;
-    else if (ABIStr.endswith("d"))
+    else if (ABIStr.ends_with("d"))
       ABIFLen = 64;
     return createRISCVTargetCodeGenInfo(CGM, XLen, ABIFLen);
   }
@@ -308,9 +308,9 @@ createTargetCodeGenInfo(CodeGenModule &CGM) {
   case llvm::Triple::loongarch64: {
     StringRef ABIStr = Target.getABI();
     unsigned ABIFRLen = 0;
-    if (ABIStr.endswith("f"))
+    if (ABIStr.ends_with("f"))
       ABIFRLen = 32;
-    else if (ABIStr.endswith("d"))
+    else if (ABIStr.ends_with("d"))
       ABIFRLen = 64;
     return createLoongArchTargetCodeGenInfo(
         CGM, Target.getPointerWidth(LangAS::Default), ABIFRLen);
@@ -1715,7 +1715,7 @@ static void AppendTargetMangling(const CodeGenModule &CGM,
   llvm::sort(Info.Features, [&Target](StringRef LHS, StringRef RHS) {
     // Multiversioning doesn't allow "no-${feature}", so we can
     // only have "+" prefixes here.
-    assert(LHS.startswith("+") && RHS.startswith("+") &&
+    assert(LHS.starts_with("+") && RHS.starts_with("+") &&
            "Features should always have a prefix.");
     return Target.multiVersionSortPriority(LHS.substr(1)) >
            Target.multiVersionSortPriority(RHS.substr(1));
@@ -1769,7 +1769,7 @@ static void AppendTargetClonesMangling(const CodeGenModule &CGM,
   } else {
     Out << '.';
     StringRef FeatureStr = Attr->getFeatureStr(VersionIndex);
-    if (FeatureStr.startswith("arch="))
+    if (FeatureStr.starts_with("arch="))
       Out << "arch_" << FeatureStr.substr(sizeof("arch=") - 1);
     else
       Out << FeatureStr;
@@ -3828,7 +3828,7 @@ namespace {
       if (!BuiltinID || !BI.isLibFunction(BuiltinID))
         return false;
       StringRef BuiltinName = BI.getName(BuiltinID);
-      if (BuiltinName.startswith("__builtin_") &&
+      if (BuiltinName.starts_with("__builtin_") &&
           Name == BuiltinName.slice(strlen("__builtin_"), StringRef::npos)) {
         return true;
       }
@@ -4164,7 +4164,7 @@ void CodeGenModule::emitMultiVersionFunctions() {
               Feature.push_back(CurFeat.trim());
           }
         } else {
-          if (Version.startswith("arch="))
+          if (Version.starts_with("arch="))
             Architecture = Version.drop_front(sizeof("arch=") - 1);
           else if (Version != "default")
             Feature.push_back(Version);

clang/lib/CodeGen/CoverageMappingGen.cpp

@@ -1623,8 +1623,12 @@ static void dump(llvm::raw_ostream &OS, StringRef FunctionName,
       OS << "Gap,";
       break;
     case CounterMappingRegion::BranchRegion:
+    case CounterMappingRegion::MCDCBranchRegion:
       OS << "Branch,";
       break;
+    case CounterMappingRegion::MCDCDecisionRegion:
+      OS << "Decision,";
+      break;
     }
 
     OS << "File " << R.FileID << ", " << R.LineStart << ":" << R.ColumnStart

clang/lib/CodeGen/Targets/SPIR.cpp

@@ -146,14 +146,14 @@ static llvm::Type *getSPIRVImageType(llvm::LLVMContext &Ctx, StringRef BaseType,
 
   // Choose the dimension of the image--this corresponds to the Dim enum in
   // SPIR-V (first integer parameter of OpTypeImage).
-  if (OpenCLName.startswith("image2d"))
+  if (OpenCLName.starts_with("image2d"))
     IntParams[0] = 1; // 1D
-  else if (OpenCLName.startswith("image3d"))
+  else if (OpenCLName.starts_with("image3d"))
     IntParams[0] = 2; // 2D
   else if (OpenCLName == "image1d_buffer")
     IntParams[0] = 5; // Buffer
   else
-    assert(OpenCLName.startswith("image1d") && "Unknown image type");
+    assert(OpenCLName.starts_with("image1d") && "Unknown image type");
 
   // Set the other integer parameters of OpTypeImage if necessary. Note that the
   // OpenCL image types don't provide any information for the Sampled or