Merge pull request #358 from AMD-Lightning-Internal/upstream_merge_202501291026

merge main into amd-staging

Author: ronlieb

clang/docs/LanguageExtensions.rst

@@ -757,6 +757,9 @@ The integer elementwise intrinsics, including ``__builtin_elementwise_popcount``
 ``__builtin_elementwise_bitreverse``, ``__builtin_elementwise_add_sat``,
 ``__builtin_elementwise_sub_sat`` can be called in a ``constexpr`` context.
 
+No implicit promotion of integer types takes place. The mixing of integer types
+of different sizes and signs is forbidden in binary and ternary builtins.
+
 ============================================== ====================================================================== =========================================
          Name                                   Operation                                                             Supported element types
 ============================================== ====================================================================== =========================================

clang/include/clang/Basic/BuiltinsAMDGPU.def

@@ -504,6 +504,8 @@ TARGET_BUILTIN(__builtin_amdgcn_global_load_tr_b128_v4i16, "V4sV4s*1", "nc", "gf
 TARGET_BUILTIN(__builtin_amdgcn_global_load_tr_b128_v4f16, "V4hV4h*1", "nc", "gfx12-insts,wavefrontsize64")
 TARGET_BUILTIN(__builtin_amdgcn_global_load_tr_b128_v4bf16, "V4yV4y*1", "nc", "gfx12-insts,wavefrontsize64")
 
+TARGET_BUILTIN(__builtin_amdgcn_ds_bpermute_fi_b32, "iii", "nc", "gfx12-insts")
+
 //===----------------------------------------------------------------------===//
 // WMMA builtins.
 // Postfix w32 indicates the builtin requires wavefront size of 32.

clang/include/clang/Sema/Sema.h

@@ -2331,7 +2331,8 @@ class Sema final : public SemaBase {
                          const FunctionProtoType *Proto);
 
   /// \param FPOnly restricts the arguments to floating-point types.
-  bool BuiltinVectorMath(CallExpr *TheCall, QualType &Res, bool FPOnly = false);
+  std::optional<QualType> BuiltinVectorMath(CallExpr *TheCall,
+                                            bool FPOnly = false);
   bool BuiltinVectorToScalarMath(CallExpr *TheCall);
 
   void checkLifetimeCaptureBy(FunctionDecl *FDecl, bool IsMemberFunction,
@@ -7499,10 +7500,15 @@ class Sema final : public SemaBase {
     return K == ConditionKind::Switch ? Context.IntTy : Context.BoolTy;
   }
 
-  // UsualUnaryConversions - promotes integers (C99 6.3.1.1p2) and converts
-  // functions and arrays to their respective pointers (C99 6.3.2.1).
+  // UsualUnaryConversions - promotes integers (C99 6.3.1.1p2), converts
+  // functions and arrays to their respective pointers (C99 6.3.2.1), and
+  // promotes floating-piont types according to the language semantics.
   ExprResult UsualUnaryConversions(Expr *E);
 
+  // UsualUnaryFPConversions - promotes floating-point types according to the
+  // current language semantics.
+  ExprResult UsualUnaryFPConversions(Expr *E);
+
   /// CallExprUnaryConversions - a special case of an unary conversion
   /// performed on a function designator of a call expression.
   ExprResult CallExprUnaryConversions(Expr *E);
@@ -7565,6 +7571,11 @@ class Sema final : public SemaBase {
   ExprResult DefaultVariadicArgumentPromotion(Expr *E, VariadicCallType CT,
                                               FunctionDecl *FDecl);
 
+  // Check that the usual arithmetic conversions can be performed on this pair
+  // of expressions that might be of enumeration type.
+  void checkEnumArithmeticConversions(Expr *LHS, Expr *RHS, SourceLocation Loc,
+                                      Sema::ArithConvKind ACK);
+
   // UsualArithmeticConversions - performs the UsualUnaryConversions on it's
   // operands and then handles various conversions that are common to binary
   // operators (C99 6.3.1.8). If both operands aren't arithmetic, this

clang/lib/AST/ByteCode/Compiler.cpp

@@ -6175,7 +6175,7 @@ bool Compiler<Emitter>::visitDeclRef(const ValueDecl *D, const Expr *E) {
     }
 
     if (D->getType()->isReferenceType())
-      return false; // FIXME: Do we need to emit InvalidDeclRef?
+      return this->emitDummyPtr(D, E);
   }
 
   // In case we need to re-visit a declaration.

clang/lib/AST/ByteCode/Descriptor.cpp

@@ -409,7 +409,8 @@ QualType Descriptor::getElemQualType() const {
   assert(isArray());
   QualType T = getType();
   if (T->isPointerOrReferenceType())
-    return T->getPointeeType();
+    T = T->getPointeeType();
+
   if (const auto *AT = T->getAsArrayTypeUnsafe()) {
     // For primitive arrays, we don't save a QualType at all,
     // just a PrimType. Try to figure out the QualType here.
@@ -424,7 +425,8 @@ QualType Descriptor::getElemQualType() const {
     return CT->getElementType();
   if (const auto *CT = T->getAs<VectorType>())
     return CT->getElementType();
-  llvm_unreachable("Array that's not an array/complex/vector type?");
+
+  return T;
 }
 
 SourceLocation Descriptor::getLocation() const {

clang/lib/AST/ByteCode/Interp.cpp

@@ -68,6 +68,9 @@ static bool diagnoseUnknownDecl(InterpState &S, CodePtr OpPC,
   const SourceInfo &E = S.Current->getSource(OpPC);
 
   if (isa<ParmVarDecl>(D)) {
+    if (D->getType()->isReferenceType())
+      return false;
+
     if (S.getLangOpts().CPlusPlus11) {
       S.FFDiag(E, diag::note_constexpr_function_param_value_unknown) << D;
       S.Note(D->getLocation(), diag::note_declared_at) << D->getSourceRange();
@@ -1287,6 +1290,12 @@ bool Call(InterpState &S, CodePtr OpPC, const Function *Func,
 
     const Pointer &ThisPtr = S.Stk.peek<Pointer>(ThisOffset);
 
+    // C++23 [expr.const]p5.6
+    // an invocation of a virtual function ([class.virtual]) for an object whose
+    // dynamic type is constexpr-unknown;
+    if (ThisPtr.isDummy() && Func->isVirtual())
+      return false;
+
     // If the current function is a lambda static invoker and
     // the function we're about to call is a lambda call operator,
     // skip the CheckInvoke, since the ThisPtr is a null pointer
@@ -1661,17 +1670,6 @@ bool GetTypeidPtr(InterpState &S, CodePtr OpPC, const Type *TypeInfoType) {
   if (!P.isBlockPointer())
     return false;
 
-  if (P.isDummy()) {
-    QualType StarThisType =
-        S.getASTContext().getLValueReferenceType(P.getType());
-    S.FFDiag(S.Current->getSource(OpPC),
-             diag::note_constexpr_polymorphic_unknown_dynamic_type)
-        << AK_TypeId
-        << P.toAPValue(S.getASTContext())
-               .getAsString(S.getASTContext(), StarThisType);
-    return false;
-  }
-
   S.Stk.push<Pointer>(P.getType().getTypePtr(), TypeInfoType);
   return true;
 }

clang/lib/AST/ByteCode/InterpBuiltin.cpp

@@ -1538,9 +1538,12 @@ static bool interp__builtin_constant_p(InterpState &S, CodePtr OpPC,
   if (ArgType->isIntegralOrEnumerationType() || ArgType->isFloatingType() ||
       ArgType->isAnyComplexType() || ArgType->isPointerType() ||
       ArgType->isNullPtrType()) {
+    auto PrevDiags = S.getEvalStatus().Diag;
+    S.getEvalStatus().Diag = nullptr;
     InterpStack Stk;
     Compiler<EvalEmitter> C(S.Ctx, S.P, S, Stk);
     auto Res = C.interpretExpr(Arg, /*ConvertResultToRValue=*/Arg->isGLValue());
+    S.getEvalStatus().Diag = PrevDiags;
     if (Res.isInvalid()) {
       C.cleanup();
       Stk.clear();

clang/lib/AST/ByteCode/Pointer.cpp

@@ -209,6 +209,10 @@ APValue Pointer::toAPValue(const ASTContext &ASTCtx) const {
     return ASTCtx.toCharUnitsFromBits(Layout.getFieldOffset(FieldIndex));
   };
 
+  bool UsePath = true;
+  if (getType()->isLValueReferenceType())
+    UsePath = false;
+
   // Build the path into the object.
   Pointer Ptr = *this;
   while (Ptr.isField() || Ptr.isArrayElement()) {
@@ -217,38 +221,42 @@ APValue Pointer::toAPValue(const ASTContext &ASTCtx) const {
       // An array root may still be an array element itself.
       if (Ptr.isArrayElement()) {
         Ptr = Ptr.expand();
+        const Descriptor *Desc = Ptr.getFieldDesc();
         unsigned Index = Ptr.getIndex();
-        Path.push_back(APValue::LValuePathEntry::ArrayIndex(Index));
-        QualType ElemType = Ptr.getFieldDesc()->getElemQualType();
+        QualType ElemType = Desc->getElemQualType();
         Offset += (Index * ASTCtx.getTypeSizeInChars(ElemType));
+        if (Ptr.getArray().getType()->isArrayType())
+          Path.push_back(APValue::LValuePathEntry::ArrayIndex(Index));
         Ptr = Ptr.getArray();
       } else {
-        Path.push_back(APValue::LValuePathEntry(
-            {Ptr.getFieldDesc()->asDecl(), /*IsVirtual=*/false}));
+        const Descriptor *Desc = Ptr.getFieldDesc();
+        const auto *Dcl = Desc->asDecl();
+        Path.push_back(APValue::LValuePathEntry({Dcl, /*IsVirtual=*/false}));
 
-        if (const auto *FD =
-                dyn_cast_if_present<FieldDecl>(Ptr.getFieldDesc()->asDecl()))
+        if (const auto *FD = dyn_cast_if_present<FieldDecl>(Dcl))
           Offset += getFieldOffset(FD);
 
         Ptr = Ptr.getBase();
       }
     } else if (Ptr.isArrayElement()) {
       Ptr = Ptr.expand();
+      const Descriptor *Desc = Ptr.getFieldDesc();
       unsigned Index;
       if (Ptr.isOnePastEnd())
         Index = Ptr.getArray().getNumElems();
       else
         Index = Ptr.getIndex();
 
-      QualType ElemType = Ptr.getFieldDesc()->getElemQualType();
+      QualType ElemType = Desc->getElemQualType();
       Offset += (Index * ASTCtx.getTypeSizeInChars(ElemType));
-      Path.push_back(APValue::LValuePathEntry::ArrayIndex(Index));
+      if (Ptr.getArray().getType()->isArrayType())
+        Path.push_back(APValue::LValuePathEntry::ArrayIndex(Index));
       Ptr = Ptr.getArray();
     } else {
+      const Descriptor *Desc = Ptr.getFieldDesc();
       bool IsVirtual = false;
 
       // Create a path entry for the field.
-      const Descriptor *Desc = Ptr.getFieldDesc();
       if (const auto *BaseOrMember = Desc->asDecl()) {
         if (const auto *FD = dyn_cast<FieldDecl>(BaseOrMember)) {
           Ptr = Ptr.getBase();
@@ -281,8 +289,11 @@ APValue Pointer::toAPValue(const ASTContext &ASTCtx) const {
   // Just invert the order of the elements.
   std::reverse(Path.begin(), Path.end());
 
-  return APValue(Base, Offset, Path, /*IsOnePastEnd=*/isOnePastEnd(),
-                 /*IsNullPtr=*/false);
+  if (UsePath)
+    return APValue(Base, Offset, Path,
+                   /*IsOnePastEnd=*/!isElementPastEnd() && isOnePastEnd());
+
+  return APValue(Base, Offset, APValue::NoLValuePath());
 }
 
 void Pointer::print(llvm::raw_ostream &OS) const {

clang/lib/AST/ByteCode/Pointer.h

@@ -630,8 +630,7 @@ class Pointer {
     if (isUnknownSizeArray())
       return false;
 
-    return isElementPastEnd() || isPastEnd() ||
-           (getSize() == getOffset() && !isZeroSizeArray());
+    return isPastEnd() || (getSize() == getOffset() && !isZeroSizeArray());
   }
 
   /// Checks if the pointer points past the end of the object.

clang/lib/Basic/Targets/ARM.cpp

@@ -608,6 +608,8 @@ bool ARMTargetInfo::handleTargetFeatures(std::vector<std::string> &Features,
       HasBTI = 1;
     } else if (Feature == "+fullbf16") {
       HasFullBFloat16 = true;
+    } else if (Feature == "+execute-only") {
+      TLSSupported = false;
     }
   }