merge main into amd-staging

Change-Id: Ib82f0e403dd6cab302c718c05124a0531a16ccd6

.github/workflows/release-binaries-all.yml

@@ -43,6 +43,7 @@ on:
       - '.github/workflows/release-binaries.yml'
       - '.github/workflows/release-binaries-setup-stage/*'
       - '.github/workflows/release-binaries-save-stage/*'
+      - 'clang/cmake/caches/Release.cmake'
 
 concurrency:
   group: ${{ github.workflow }}-${{ github.event.pull_request.number || 'dispatch' }}

clang/docs/ReleaseNotes.rst

@@ -103,6 +103,7 @@ ABI Changes in This Version
 ---------------------------
 
 - Fixed Microsoft name mangling of placeholder, auto and decltype(auto), return types for MSVC 1920+. This change resolves incompatibilities with code compiled by MSVC 1920+ but will introduce incompatibilities with code compiled by earlier versions of Clang unless such code is built with the compiler option -fms-compatibility-version=19.14 to imitate the MSVC 1914 mangling behavior.
+- Fixed the Itanium mangling of the construction vtable name. This change will introduce incompatibilities with code compiled by Clang 19 and earlier versions, unless the -fclang-abi-compat=19 option is used. (#GH108015)
 
 AST Dumping Potentially Breaking Changes
 ----------------------------------------

clang/include/clang/Basic/LangOptions.h

@@ -239,6 +239,12 @@ class LangOptionsBase {
     /// in the initializers of members of local classes.
     Ver18,
 
+    /// Attempt to be ABI-compatible with code generated by Clang 19.0.x.
+    /// This causes clang to:
+    ///   - Incorrectly mangles the 'base type' substitutions of the CXX
+    ///   construction vtable because it hasn't added 'type' as a substitution.
+    Ver19,
+
     /// Conform to the underlying platform's C and C++ ABIs as closely
     /// as we can.
     Latest

clang/lib/AST/ByteCode/Compiler.cpp

@@ -697,6 +697,14 @@ bool Compiler<Emitter>::VisitCastExpr(const CastExpr *CE) {
     const auto *TargetSemantics = &Ctx.getFloatSemantics(CE->getType());
     return this->emitCastFixedPointFloating(TargetSemantics, CE);
   }
+  case CK_FixedPointCast: {
+    if (!this->visit(SubExpr))
+      return false;
+    auto Sem = Ctx.getASTContext().getFixedPointSemantics(CE->getType());
+    uint32_t I;
+    std::memcpy(&I, &Sem, sizeof(Sem));
+    return this->emitCastFixedPoint(I, CE);
+  }
 
   case CK_ToVoid:
     return discard(SubExpr);
@@ -1494,25 +1502,41 @@ bool Compiler<Emitter>::VisitFixedPointBinOp(const BinaryOperator *E) {
   assert(LHS->getType()->isFixedPointType() ||
          RHS->getType()->isFixedPointType());
 
+  auto LHSSema = Ctx.getASTContext().getFixedPointSemantics(LHS->getType());
+  auto RHSSema = Ctx.getASTContext().getFixedPointSemantics(RHS->getType());
+
   if (!this->visit(LHS))
     return false;
   if (!LHS->getType()->isFixedPointType()) {
-    auto Sem = Ctx.getASTContext().getFixedPointSemantics(LHS->getType());
     uint32_t I;
-    std::memcpy(&I, &Sem, sizeof(Sem));
+    std::memcpy(&I, &LHSSema, sizeof(llvm::FixedPointSemantics));
     if (!this->emitCastIntegralFixedPoint(classifyPrim(LHS->getType()), I, E))
       return false;
   }
+
   if (!this->visit(RHS))
     return false;
   if (!RHS->getType()->isFixedPointType()) {
-    auto Sem = Ctx.getASTContext().getFixedPointSemantics(RHS->getType());
     uint32_t I;
-    std::memcpy(&I, &Sem, sizeof(Sem));
+    std::memcpy(&I, &RHSSema, sizeof(llvm::FixedPointSemantics));
     if (!this->emitCastIntegralFixedPoint(classifyPrim(RHS->getType()), I, E))
       return false;
   }
 
+  // Convert the result to the target semantics.
+  auto ConvertResult = [&](bool R) -> bool {
+    if (!R)
+      return false;
+    auto ResultSema = Ctx.getASTContext().getFixedPointSemantics(E->getType());
+    auto CommonSema = LHSSema.getCommonSemantics(RHSSema);
+    if (ResultSema != CommonSema) {
+      uint32_t I;
+      std::memcpy(&I, &ResultSema, sizeof(ResultSema));
+      return this->emitCastFixedPoint(I, E);
+    }
+    return true;
+  };
+
   switch (E->getOpcode()) {
   case BO_EQ:
     return this->emitEQFixedPoint(E);
@@ -1528,6 +1552,9 @@ bool Compiler<Emitter>::VisitFixedPointBinOp(const BinaryOperator *E) {
   case BO_GE:
     return this->emitGEFixedPoint(E);
 #endif
+  case BO_Add:
+    return ConvertResult(this->emitAddFixedPoint(E));
+
   default:
     return this->emitInvalid(E);
   }

clang/lib/AST/ByteCode/FixedPoint.h

@@ -47,15 +47,33 @@ class FixedPoint final {
   void print(llvm::raw_ostream &OS) const { OS << V; }
 
   APValue toAPValue(const ASTContext &) const { return APValue(V); }
-  APSInt toAPSInt(unsigned BitWidth) const { return V.getValue(); }
+  APSInt toAPSInt(unsigned BitWidth = 0) const { return V.getValue(); }
 
   unsigned bitWidth() const { return V.getWidth(); }
   bool isSigned() const { return V.isSigned(); }
+  bool isZero() const { return V.getValue().isZero(); }
+  bool isNegative() const { return V.getValue().isNegative(); }
+  bool isPositive() const { return V.getValue().isNonNegative(); }
+  bool isMin() const {
+    return V.getValue() == APSInt::getMinValue(V.getSemantics().getWidth(),
+                                               !V.getSemantics().isSigned());
+  }
+
+  FixedPoint truncate(unsigned BitWidth) const { return *this; }
+
+  FixedPoint toSemantics(const llvm::FixedPointSemantics &Sem,
+                         bool *Overflow) const {
+    return FixedPoint(V.convert(Sem, Overflow));
+  }
 
   llvm::APFloat toFloat(const llvm::fltSemantics *Sem) const {
     return V.convertToFloat(*Sem);
   }
 
+  std::string toDiagnosticString(const ASTContext &Ctx) const {
+    return V.toString();
+  }
+
   ComparisonCategoryResult compare(const FixedPoint &Other) const {
     if (Other.V == V)
       return ComparisonCategoryResult::Equal;
@@ -67,6 +85,27 @@ class FixedPoint final {
     *R = FixedPoint(A.V.negate(&Overflow));
     return Overflow;
   }
+
+  static bool add(const FixedPoint A, const FixedPoint B, unsigned Bits,
+                  FixedPoint *R) {
+    bool Overflow = false;
+    *R = FixedPoint(A.V.add(B.V, &Overflow));
+    return Overflow;
+  }
+  static bool sub(const FixedPoint A, const FixedPoint B, unsigned Bits,
+                  FixedPoint *R) {
+    return true;
+  }
+  static bool mul(const FixedPoint A, const FixedPoint B, unsigned Bits,
+                  FixedPoint *R) {
+    return true;
+  }
+  static bool div(const FixedPoint A, const FixedPoint B, unsigned Bits,
+                  FixedPoint *R) {
+    return true;
+  }
+  static bool increment(const FixedPoint &A, FixedPoint *R) { return true; }
+  static bool decrement(const FixedPoint &A, FixedPoint *R) { return true; }
 };
 
 inline FixedPoint getSwappedBytes(FixedPoint F) { return F; }

clang/lib/AST/ByteCode/Interp.h

@@ -2161,6 +2161,32 @@ inline bool CastFP(InterpState &S, CodePtr OpPC, const llvm::fltSemantics *Sem,
   return true;
 }
 
+inline bool CastFixedPoint(InterpState &S, CodePtr OpPC, uint32_t FPS) {
+  FixedPointSemantics TargetSemantics(0, 0, false, false, false);
+  std::memcpy(&TargetSemantics, &FPS, sizeof(TargetSemantics));
+
+  const auto &Source = S.Stk.pop<FixedPoint>();
+
+  bool Overflow;
+  FixedPoint Result = Source.toSemantics(TargetSemantics, &Overflow);
+
+  if (Overflow) {
+    const Expr *E = S.Current->getExpr(OpPC);
+    if (S.checkingForUndefinedBehavior()) {
+      S.getASTContext().getDiagnostics().Report(
+          E->getExprLoc(), diag::warn_fixedpoint_constant_overflow)
+          << Result.toDiagnosticString(S.getASTContext()) << E->getType();
+    }
+    S.CCEDiag(E, diag::note_constexpr_overflow)
+        << Result.toDiagnosticString(S.getASTContext()) << E->getType();
+    if (!S.noteUndefinedBehavior())
+      return false;
+  }
+
+  S.Stk.push<FixedPoint>(Result);
+  return true;
+}
+
 /// Like Cast(), but we cast to an arbitrary-bitwidth integral, so we need
 /// to know what bitwidth the result should be.
 template <PrimType Name, class T = typename PrimConv<Name>::T>

clang/lib/AST/ByteCode/Opcodes.td

@@ -98,7 +98,7 @@ def FloatTypeClass : TypeClass {
 }
 
 def AluTypeClass : TypeClass {
-  let Types = !listconcat(IntegerTypeClass.Types, [Bool]);
+  let Types = !listconcat(IntegerTypeClass.Types, [Bool], [FixedPoint]);
 }
 
 def PtrTypeClass : TypeClass {
@@ -110,7 +110,7 @@ def NonPtrTypeClass : TypeClass {
 }
 
 def AllTypeClass : TypeClass {
-  let Types = !listconcat(AluTypeClass.Types, PtrTypeClass.Types, FloatTypeClass.Types, [FixedPoint]);
+  let Types = !listconcat(AluTypeClass.Types, PtrTypeClass.Types, FloatTypeClass.Types);
 }
 
 def ComparableTypeClass : TypeClass {
@@ -626,6 +626,10 @@ def CastFP : Opcode {
   let Args = [ArgFltSemantics, ArgRoundingMode];
 }
 
+def CastFixedPoint : Opcode {
+  let Args = [ArgUint32];
+}
+
 def FixedSizeIntegralTypes : TypeClass {
   let Types = [Uint8, Sint8, Uint16, Sint16, Uint32, Sint32, Uint64, Sint64, Bool];
 }

clang/lib/AST/ItaniumMangle.cpp

@@ -464,7 +464,7 @@ class CXXNameMangler {
   void mangleSeqID(unsigned SeqID);
   void mangleName(GlobalDecl GD);
   void mangleType(QualType T);
-  void mangleNameOrStandardSubstitution(const NamedDecl *ND);
+  void mangleCXXRecordDecl(const CXXRecordDecl *Record);
   void mangleLambdaSig(const CXXRecordDecl *Lambda);
   void mangleModuleNamePrefix(StringRef Name, bool IsPartition = false);
   void mangleVendorQualifier(StringRef Name);
@@ -3029,9 +3029,13 @@ void CXXNameMangler::mangleType(QualType T) {
     addSubstitution(T);
 }
 
-void CXXNameMangler::mangleNameOrStandardSubstitution(const NamedDecl *ND) {
-  if (!mangleStandardSubstitution(ND))
-    mangleName(ND);
+void CXXNameMangler::mangleCXXRecordDecl(const CXXRecordDecl *Record) {
+  if (mangleSubstitution(Record))
+    return;
+  mangleName(Record);
+  if (isCompatibleWith(LangOptions::ClangABI::Ver19))
+    return;
+  addSubstitution(Record);
 }
 
 void CXXNameMangler::mangleType(const BuiltinType *T) {
@@ -7309,15 +7313,15 @@ void ItaniumMangleContextImpl::mangleCXXVTable(const CXXRecordDecl *RD,
   // <special-name> ::= TV <type>  # virtual table
   CXXNameMangler Mangler(*this, Out);
   Mangler.getStream() << "_ZTV";
-  Mangler.mangleNameOrStandardSubstitution(RD);
+  Mangler.mangleCXXRecordDecl(RD);
 }
 
 void ItaniumMangleContextImpl::mangleCXXVTT(const CXXRecordDecl *RD,
                                             raw_ostream &Out) {
   // <special-name> ::= TT <type>  # VTT structure
   CXXNameMangler Mangler(*this, Out);
   Mangler.getStream() << "_ZTT";
-  Mangler.mangleNameOrStandardSubstitution(RD);
+  Mangler.mangleCXXRecordDecl(RD);
 }
 
 void ItaniumMangleContextImpl::mangleCXXCtorVTable(const CXXRecordDecl *RD,
@@ -7327,10 +7331,10 @@ void ItaniumMangleContextImpl::mangleCXXCtorVTable(const CXXRecordDecl *RD,
   // <special-name> ::= TC <type> <offset number> _ <base type>
   CXXNameMangler Mangler(*this, Out);
   Mangler.getStream() << "_ZTC";
-  Mangler.mangleNameOrStandardSubstitution(RD);
+  Mangler.mangleCXXRecordDecl(RD);
   Mangler.getStream() << Offset;
   Mangler.getStream() << '_';
-  Mangler.mangleNameOrStandardSubstitution(Type);
+  Mangler.mangleCXXRecordDecl(Type);
 }
 
 void ItaniumMangleContextImpl::mangleCXXRTTI(QualType Ty, raw_ostream &Out) {

clang/lib/CodeGen/CodeGenModule.cpp

@@ -5633,8 +5633,9 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D,
     emitter->finalize(GV);
 
   // If it is safe to mark the global 'constant', do so now.
-  GV->setConstant(!NeedsGlobalCtor && !NeedsGlobalDtor &&
-                  D->getType().isConstantStorage(getContext(), true, true));
+  GV->setConstant((D->hasAttr<CUDAConstantAttr>() && LangOpts.CUDAIsDevice) ||
+                  (!NeedsGlobalCtor && !NeedsGlobalDtor &&
+                   D->getType().isConstantStorage(getContext(), true, true)));
 
   // If it is in a read-only section, mark it 'constant'.
   if (const SectionAttr *SA = D->getAttr<SectionAttr>()) {

clang/lib/Frontend/CompilerInvocation.cpp

@@ -3891,6 +3891,9 @@ void CompilerInvocationBase::GenerateLangArgs(const LangOptions &Opts,
   case LangOptions::ClangABI::Ver18:
     GenerateArg(Consumer, OPT_fclang_abi_compat_EQ, "18.0");
     break;
+  case LangOptions::ClangABI::Ver19:
+    GenerateArg(Consumer, OPT_fclang_abi_compat_EQ, "19.0");
+    break;
   case LangOptions::ClangABI::Latest:
     break;
   }
@@ -4482,6 +4485,8 @@ bool CompilerInvocation::ParseLangArgs(LangOptions &Opts, ArgList &Args,
         Opts.setClangABICompat(LangOptions::ClangABI::Ver17);
       else if (Major <= 18)
         Opts.setClangABICompat(LangOptions::ClangABI::Ver18);
+      else if (Major <= 19)
+        Opts.setClangABICompat(LangOptions::ClangABI::Ver19);
     } else if (Ver != "latest") {
       Diags.Report(diag::err_drv_invalid_value)
           << A->getAsString(Args) << A->getValue();

clang/lib/Serialization/ASTWriterDecl.cpp

@@ -1785,6 +1785,18 @@ void ASTDeclWriter::VisitClassTemplateSpecializationDecl(
   if (ArgsWritten)
     Record.AddASTTemplateArgumentListInfo(ArgsWritten);
 
+  // Mention the implicitly generated C++ deduction guide to make sure the
+  // deduction guide will be rewritten as expected.
+  //
+  // FIXME: Would it be more efficient to add a callback register function
+  // in sema to register the deduction guide?
+  if (Writer.isWritingStdCXXNamedModules()) {
+    auto Name = Context.DeclarationNames.getCXXDeductionGuideName(
+        D->getSpecializedTemplate());
+    for (auto *DG : D->getDeclContext()->noload_lookup(Name))
+      Writer.GetDeclRef(DG->getCanonicalDecl());
+  }
+
   Code = serialization::DECL_CLASS_TEMPLATE_SPECIALIZATION;
 }
 

clang/test/AST/ByteCode/fixed-point.cpp

@@ -35,3 +35,19 @@ namespace FloatToFixedPointCast {
   constexpr float sf2f = sf2;
   static_assert(sf2f == 0.5);
 }
+
+namespace BinOps {
+  constexpr _Accum A = 13;
+  static_assert(A + 1 == 14.0k);
+  static_assert(1 + A == 14.0k);
+  static_assert((A + A) == 26);
+
+  static_assert(A + 100000 == 14.0k); // both-error {{is not an integral constant expression}} \
+                                      // both-note {{is outside the range of representable values}}
+}
+
+namespace FixedPointCasts {
+  constexpr _Fract B = 0.3;
+  constexpr _Accum A = B;
+  constexpr _Fract C = A;
+}

clang/test/CodeGenCUDA/address-spaces.cu

@@ -9,7 +9,7 @@
 // CHECK: @i ={{.*}} addrspace(1) externally_initialized global
 __device__ int i;
 
-// CHECK: @j ={{.*}} addrspace(4) externally_initialized global
+// CHECK: @j ={{.*}} addrspace(4) externally_initialized constant
 __constant__ int j;
 
 // CHECK: @k ={{.*}} addrspace(3) global

clang/test/CodeGenCUDA/amdgpu-visibility.cu

@@ -4,11 +4,11 @@
 
 #include "Inputs/cuda.h"
 
-// CHECK-DEFAULT: @c ={{.*}} addrspace(4) externally_initialized global
+// CHECK-DEFAULT: @c ={{.*}} addrspace(4) externally_initialized constant
 // CHECK-DEFAULT: @g ={{.*}} addrspace(1) externally_initialized global
-// CHECK-PROTECTED: @c = protected addrspace(4) externally_initialized global
+// CHECK-PROTECTED: @c = protected addrspace(4) externally_initialized constant
 // CHECK-PROTECTED: @g = protected addrspace(1) externally_initialized global
-// CHECK-HIDDEN: @c = protected addrspace(4) externally_initialized global
+// CHECK-HIDDEN: @c = protected addrspace(4) externally_initialized constant
 // CHECK-HIDDEN: @g = protected addrspace(1) externally_initialized global
 __constant__ int c;
 __device__ int g;

clang/test/CodeGenCUDA/anon-ns.cu

@@ -28,13 +28,13 @@
 // HIP-DAG: define weak_odr {{.*}}void @[[KTX:_Z2ktIN12_GLOBAL__N_11XEEvT_\.intern\.b04fd23c98500190]](
 // HIP-DAG: define weak_odr {{.*}}void @[[KTL:_Z2ktIN12_GLOBAL__N_1UlvE_EEvT_\.intern\.b04fd23c98500190]](
 // HIP-DAG: @[[VM:_ZN12_GLOBAL__N_12vmE\.static\.b04fd23c98500190]] = addrspace(1) externally_initialized global
-// HIP-DAG: @[[VC:_ZN12_GLOBAL__N_12vcE\.static\.b04fd23c98500190]] = addrspace(4) externally_initialized global
+// HIP-DAG: @[[VC:_ZN12_GLOBAL__N_12vcE\.static\.b04fd23c98500190]] = addrspace(4) externally_initialized constant
 // HIP-DAG: @[[VT:_Z2vtIN12_GLOBAL__N_11XEE\.static\.b04fd23c98500190]] = addrspace(1) externally_initialized global
 
 // CUDA-DAG: define weak_odr {{.*}}void @[[KERN:_ZN12_GLOBAL__N_16kernelEv__intern__b04fd23c98500190]](
 // CUDA-DAG: define weak_odr {{.*}}void @[[KTX:_Z2ktIN12_GLOBAL__N_11XEEvT___intern__b04fd23c98500190]](
 // CUDA-DAG: define weak_odr {{.*}}void @[[KTL:_Z2ktIN12_GLOBAL__N_1UlvE_EEvT___intern__b04fd23c98500190]](
-// CUDA-DAG: @[[VC:_ZN12_GLOBAL__N_12vcE__static__b04fd23c98500190]] = addrspace(4) externally_initialized global
+// CUDA-DAG: @[[VC:_ZN12_GLOBAL__N_12vcE__static__b04fd23c98500190]] = addrspace(4) externally_initialized constant
 // CUDA-DAG: @[[VT:_Z2vtIN12_GLOBAL__N_11XEE__static__b04fd23c98500190]] = addrspace(1) externally_initialized global
 
 // COMMON-DAG: @_ZN12_GLOBAL__N_12vdE = internal addrspace(1) global