Merge from 'main' to 'sycl-web' (85 commits)

Author: iclsrc

.github/workflows/libc-fullbuild-tests.yml

@@ -21,37 +21,37 @@ jobs:
           - os: ubuntu-24.04
             build_type: Debug
             ccache-variant: sccache
-            c_compiler: clang-20
-            cpp_compiler: clang++-20
+            c_compiler: clang-22
+            cpp_compiler: clang++-22
             target: x86_64-unknown-linux-llvm
             include_scudo: ON
           - os: ubuntu-24.04
             build_type: Release
             ccache-variant: sccache
-            c_compiler: clang-20
-            cpp_compiler: clang++-20
+            c_compiler: clang-22
+            cpp_compiler: clang++-22
             target: x86_64-unknown-linux-llvm
             include_scudo: ON
           - os: ubuntu-24.04
             build_type: MinSizeRel
             ccache-variant: sccache
-            c_compiler: clang-20
-            cpp_compiler: clang++-20
+            c_compiler: clang-22
+            cpp_compiler: clang++-22
             target: x86_64-unknown-linux-llvm
             include_scudo: ON
           # TODO: remove ccache logic when https://github.com/hendrikmuhs/ccache-action/issues/279 is resolved.
           - os: ubuntu-24.04-arm
             build_type: Debug
             ccache-variant: ccache
-            c_compiler: clang-20
-            cpp_compiler: clang++-20
+            c_compiler: clang-22
+            cpp_compiler: clang++-22
             target: aarch64-unknown-linux-llvm
             include_scudo: ON
           - os: ubuntu-24.04
             build_type: Debug
             ccache-variant: ccache
-            c_compiler: clang-20
-            cpp_compiler: clang++-20
+            c_compiler: clang-22
+            cpp_compiler: clang++-22
             target: x86_64-unknown-uefi-llvm
             include_scudo: OFF
           # TODO: add back gcc build when it is fixed
@@ -81,7 +81,7 @@ jobs:
       run: |
         wget https://apt.llvm.org/llvm.sh
         chmod +x llvm.sh
-        sudo ./llvm.sh 20
+        sudo ./llvm.sh 22
         sudo apt-get update
         sudo apt-get install -y libmpfr-dev libgmp-dev libmpc-dev ninja-build linux-libc-dev
         sudo ln -sf /usr/include/$(uname -p)-linux-gnu/asm /usr/include/asm

clang/include/clang/AST/OpenACCClause.h

@@ -837,23 +837,43 @@ class OpenACCClauseWithVarList : public OpenACCClauseWithExprs {
 
 class OpenACCPrivateClause final
     : public OpenACCClauseWithVarList,
-      private llvm::TrailingObjects<OpenACCPrivateClause, Expr *> {
+      private llvm::TrailingObjects<OpenACCPrivateClause, Expr *, VarDecl *> {
   friend TrailingObjects;
 
   OpenACCPrivateClause(SourceLocation BeginLoc, SourceLocation LParenLoc,
-                       ArrayRef<Expr *> VarList, SourceLocation EndLoc)
+                       ArrayRef<Expr *> VarList,
+                       ArrayRef<VarDecl *> InitRecipes, SourceLocation EndLoc)
       : OpenACCClauseWithVarList(OpenACCClauseKind::Private, BeginLoc,
                                  LParenLoc, EndLoc) {
-    setExprs(getTrailingObjects(VarList.size()), VarList);
+    assert(VarList.size() == InitRecipes.size());
+    setExprs(getTrailingObjects<Expr *>(VarList.size()), VarList);
+    llvm::uninitialized_copy(InitRecipes, getTrailingObjects<VarDecl *>());
   }
 
 public:
   static bool classof(const OpenACCClause *C) {
     return C->getClauseKind() == OpenACCClauseKind::Private;
   }
+  // Gets a list of 'made up' `VarDecl` objects that can be used by codegen to
+  // ensure that we properly initialize each of these variables.
+  ArrayRef<VarDecl *> getInitRecipes() {
+    return ArrayRef<VarDecl *>{getTrailingObjects<VarDecl *>(),
+                               getExprs().size()};
+  }
+
+  ArrayRef<VarDecl *> getInitRecipes() const {
+    return ArrayRef<VarDecl *>{getTrailingObjects<VarDecl *>(),
+                               getExprs().size()};
+  }
+
   static OpenACCPrivateClause *
   Create(const ASTContext &C, SourceLocation BeginLoc, SourceLocation LParenLoc,
-         ArrayRef<Expr *> VarList, SourceLocation EndLoc);
+         ArrayRef<Expr *> VarList, ArrayRef<VarDecl *> InitRecipes,
+         SourceLocation EndLoc);
+
+  size_t numTrailingObjects(OverloadToken<Expr *>) const {
+    return getExprs().size();
+  }
 };
 
 class OpenACCFirstPrivateClause final

clang/include/clang/Basic/BuiltinsAMDGPU.def

@@ -795,6 +795,8 @@ TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x128_fp8_fp8, "V8fV16iV16iIsV8fIbI
 TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x128_fp8_bf8, "V8fV16iV16iIsV8fIbIb", "nc", "gfx1250-insts,wavefrontsize32")
 TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x128_bf8_fp8, "V8fV16iV16iIsV8fIbIb", "nc", "gfx1250-insts,wavefrontsize32")
 TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x128_bf8_bf8, "V8fV16iV16iIsV8fIbIb", "nc", "gfx1250-insts,wavefrontsize32")
+TARGET_BUILTIN(__builtin_amdgcn_wmma_scale_f32_16x16x128_f8f6f4, "V8fIiV16iIiV16iIsV8fIiIiiIiIiiIbIb", "nc", "gfx1250-insts,wavefrontsize32")
+TARGET_BUILTIN(__builtin_amdgcn_wmma_scale16_f32_16x16x128_f8f6f4, "V8fIiV16iIiV16iIsV8fIiIiLiIiIiLiIbIb", "nc", "gfx1250-insts,wavefrontsize32")
 TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_16x16x32_f16, "V8fIbV16hIbV16hIsV8fIbIb", "nc", "gfx1250-insts,wavefrontsize32")
 TARGET_BUILTIN(__builtin_amdgcn_wmma_f16_16x16x32_f16, "V8hIbV16hIbV16hIsV8hIbIb", "nc", "gfx1250-insts,wavefrontsize32")
 TARGET_BUILTIN(__builtin_amdgcn_wmma_f32_32x16x128_f4, "V16fV16iV8iIsV16f", "nc", "gfx1250-insts,wavefrontsize32")

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -13851,7 +13851,8 @@ def err_acc_device_type_multiple_archs
 def warn_acc_var_referenced_lacks_op
     : Warning<"variable of type %0 referenced in OpenACC '%1' clause does not "
               "have a %enum_select<AccVarReferencedReason>{%DefCtor{default "
-              "constructor}|%Dtor{destructor}}2; reference has no effect">,
+              "constructor}|%CopyCtor{copy constructor}|%Dtor{destructor}}2; "
+              "reference has no effect">,
       InGroup<DiagGroup<"openacc-var-lacks-operation">>,
       DefaultError;
 

clang/include/clang/Sema/SemaOpenACC.h

@@ -176,10 +176,6 @@ class SemaOpenACC : public SemaBase {
 
     void checkFor();
 
-    //  void checkRangeFor(); ?? ERICH
-    //  const ValueDecl *checkInit();
-    //  void checkCond(const ValueDecl *Init);
-    //  void checkInc(const ValueDecl *Init);
   public:
     // Checking for non-instantiation version of a Range-for.
     ForStmtBeginChecker(SemaOpenACC &SemaRef, SourceLocation ForLoc,
@@ -241,6 +237,10 @@ class SemaOpenACC : public SemaBase {
                                 SourceLocation ClauseLoc,
                                 ArrayRef<const OpenACCClause *> Clauses);
 
+  // Creates a VarDecl with a proper default init for the purposes of a
+  // `private` clause, so it can be used to generate a recipe later.
+  VarDecl *CreateInitRecipe(const Expr *VarExpr);
+
 public:
   ComputeConstructInfo &getActiveComputeConstructInfo() {
     return ActiveComputeConstructInfo;

clang/lib/AST/ByteCode/ByteCodeEmitter.cpp

@@ -135,8 +135,8 @@ int32_t ByteCodeEmitter::getOffset(LabelTy Label) {
 /// Helper to write bytecode and bail out if 32-bit offsets become invalid.
 /// Pointers will be automatically marshalled as 32-bit IDs.
 template <typename T>
-static void emit(Program &P, std::vector<std::byte> &Code, const T &Val,
-                 bool &Success) {
+static void emit(Program &P, llvm::SmallVectorImpl<std::byte> &Code,
+                 const T &Val, bool &Success) {
   size_t ValPos = Code.size();
   size_t Size;
 
@@ -153,7 +153,7 @@ static void emit(Program &P, std::vector<std::byte> &Code, const T &Val,
   // Access must be aligned!
   assert(aligned(ValPos));
   assert(aligned(ValPos + Size));
-  Code.resize(ValPos + Size);
+  Code.resize_for_overwrite(ValPos + Size);
 
   if constexpr (!std::is_pointer_v<T>) {
     new (Code.data() + ValPos) T(Val);
@@ -166,7 +166,7 @@ static void emit(Program &P, std::vector<std::byte> &Code, const T &Val,
 /// Emits a serializable value. These usually (potentially) contain
 /// heap-allocated memory and aren't trivially copyable.
 template <typename T>
-static void emitSerialized(std::vector<std::byte> &Code, const T &Val,
+static void emitSerialized(llvm::SmallVectorImpl<std::byte> &Code, const T &Val,
                            bool &Success) {
   size_t ValPos = Code.size();
   size_t Size = align(Val.bytesToSerialize());
@@ -179,32 +179,32 @@ static void emitSerialized(std::vector<std::byte> &Code, const T &Val,
   // Access must be aligned!
   assert(aligned(ValPos));
   assert(aligned(ValPos + Size));
-  Code.resize(ValPos + Size);
+  Code.resize_for_overwrite(ValPos + Size);
 
   Val.serialize(Code.data() + ValPos);
 }
 
 template <>
-void emit(Program &P, std::vector<std::byte> &Code, const Floating &Val,
-          bool &Success) {
+void emit(Program &P, llvm::SmallVectorImpl<std::byte> &Code,
+          const Floating &Val, bool &Success) {
   emitSerialized(Code, Val, Success);
 }
 
 template <>
-void emit(Program &P, std::vector<std::byte> &Code,
+void emit(Program &P, llvm::SmallVectorImpl<std::byte> &Code,
           const IntegralAP<false> &Val, bool &Success) {
   emitSerialized(Code, Val, Success);
 }
 
 template <>
-void emit(Program &P, std::vector<std::byte> &Code, const IntegralAP<true> &Val,
-          bool &Success) {
+void emit(Program &P, llvm::SmallVectorImpl<std::byte> &Code,
+          const IntegralAP<true> &Val, bool &Success) {
   emitSerialized(Code, Val, Success);
 }
 
 template <>
-void emit(Program &P, std::vector<std::byte> &Code, const FixedPoint &Val,
-          bool &Success) {
+void emit(Program &P, llvm::SmallVectorImpl<std::byte> &Code,
+          const FixedPoint &Val, bool &Success) {
   emitSerialized(Code, Val, Success);
 }
 

clang/lib/AST/ByteCode/ByteCodeEmitter.h

@@ -88,7 +88,7 @@ class ByteCodeEmitter {
   /// Location of label relocations.
   llvm::DenseMap<LabelTy, llvm::SmallVector<unsigned, 5>> LabelRelocs;
   /// Program code.
-  std::vector<std::byte> Code;
+  llvm::SmallVector<std::byte> Code;
   /// Opcode to expression mapping.
   SourceMap SrcMap;
 

clang/lib/AST/ByteCode/Compiler.cpp

@@ -2038,7 +2038,12 @@ bool Compiler<Emitter>::visitCallArgs(ArrayRef<const Expr *> Args,
                                       const FunctionDecl *FuncDecl,
                                       bool Activate) {
   assert(VarScope->getKind() == ScopeKind::Call);
-  llvm::BitVector NonNullArgs = collectNonNullArgs(FuncDecl, Args);
+  bool HasNonNullAttr = false;
+  llvm::BitVector NonNullArgs;
+  if (FuncDecl && FuncDecl->hasAttr<NonNullAttr>()) {
+    HasNonNullAttr = true;
+    NonNullArgs = collectNonNullArgs(FuncDecl, Args);
+  }
 
   unsigned ArgIndex = 0;
   for (const Expr *Arg : Args) {
@@ -2064,7 +2069,7 @@ bool Compiler<Emitter>::visitCallArgs(ArrayRef<const Expr *> Args,
         return false;
     }
 
-    if (FuncDecl && NonNullArgs[ArgIndex]) {
+    if (HasNonNullAttr && NonNullArgs[ArgIndex]) {
       PrimType ArgT = classify(Arg).value_or(PT_Ptr);
       if (ArgT == PT_Ptr) {
         if (!this->emitCheckNonNullArg(ArgT, Arg))
@@ -5914,7 +5919,7 @@ bool Compiler<Emitter>::emitLambdaStaticInvokerBody(const CXXMethodDecl *MD) {
 
   const CXXRecordDecl *ClosureClass = MD->getParent();
   const CXXMethodDecl *LambdaCallOp = ClosureClass->getLambdaCallOperator();
-  assert(ClosureClass->captures_begin() == ClosureClass->captures_end());
+  assert(ClosureClass->captures().empty());
   const Function *Func = this->getFunction(LambdaCallOp);
   if (!Func)
     return false;

clang/lib/AST/ByteCode/Context.cpp

@@ -45,12 +45,12 @@ bool Context::isPotentialConstantExpr(State &Parent, const FunctionDecl *FD) {
   Compiler<ByteCodeEmitter>(*this, *P).compileFunc(
       FD, const_cast<Function *>(Func));
 
-  ++EvalID;
-  // And run it.
-  if (!Run(Parent, Func))
+  if (!Func->isValid())
     return false;
 
-  return Func->isValid();
+  ++EvalID;
+  // And run it.
+  return Run(Parent, Func);
 }
 
 void Context::isPotentialConstantExprUnevaluated(State &Parent, const Expr *E,
@@ -474,7 +474,7 @@ const Function *Context::getOrCreateFunction(const FunctionDecl *FuncDecl) {
     IsLambdaStaticInvoker = true;
 
     const CXXRecordDecl *ClosureClass = MD->getParent();
-    assert(ClosureClass->captures_begin() == ClosureClass->captures_end());
+    assert(ClosureClass->captures().empty());
     if (ClosureClass->isGenericLambda()) {
       const CXXMethodDecl *LambdaCallOp = ClosureClass->getLambdaCallOperator();
       assert(MD->isFunctionTemplateSpecialization() &&

clang/lib/AST/ByteCode/EvalEmitter.cpp

@@ -233,6 +233,14 @@ template <> bool EvalEmitter::emitRet<PT_Ptr>(const SourceInfo &Info) {
       return false;
     }
   } else {
+    // If this is pointing to a local variable, just return
+    // the result, even if the pointer is dead.
+    // This will later be diagnosed by CheckLValueConstantExpression.
+    if (Ptr.isBlockPointer() && !Ptr.block()->isStatic()) {
+      EvalResult.setValue(Ptr.toAPValue(Ctx.getASTContext()));
+      return true;
+    }
+
     if (!Ptr.isLive() && !Ptr.isTemporary())
       return false;