merge main into amd-staging

Change-Id: I3b824c955fe38db9e44eb4b816001e4cd91917b1

bolt/lib/Core/BinaryEmitter.cpp

@@ -46,17 +46,13 @@ BreakFunctionNames("break-funcs",
   cl::Hidden,
   cl::cat(BoltCategory));
 
-cl::list<std::string>
-    FunctionPadSpec("pad-funcs", cl::CommaSeparated,
-                    cl::desc("list of functions to pad with amount of bytes"),
-                    cl::value_desc("func1:pad1,func2:pad2,func3:pad3,..."),
-                    cl::Hidden, cl::cat(BoltCategory));
-
-cl::list<std::string> FunctionPadBeforeSpec(
-    "pad-funcs-before", cl::CommaSeparated,
-    cl::desc("list of functions to pad with amount of bytes"),
-    cl::value_desc("func1:pad1,func2:pad2,func3:pad3,..."), cl::Hidden,
-    cl::cat(BoltCategory));
+static cl::list<std::string>
+FunctionPadSpec("pad-funcs",
+  cl::CommaSeparated,
+  cl::desc("list of functions to pad with amount of bytes"),
+  cl::value_desc("func1:pad1,func2:pad2,func3:pad3,..."),
+  cl::Hidden,
+  cl::cat(BoltCategory));
 
 static cl::opt<bool> MarkFuncs(
     "mark-funcs",
@@ -74,12 +70,11 @@ X86AlignBranchBoundaryHotOnly("x86-align-branch-boundary-hot-only",
   cl::init(true),
   cl::cat(BoltOptCategory));
 
-size_t padFunction(const cl::list<std::string> &Spec,
-                   const BinaryFunction &Function) {
+size_t padFunction(const BinaryFunction &Function) {
   static std::map<std::string, size_t> FunctionPadding;
 
-  if (FunctionPadding.empty() && !Spec.empty()) {
-    for (const std::string &Spec : Spec) {
+  if (FunctionPadding.empty() && !FunctionPadSpec.empty()) {
+    for (std::string &Spec : FunctionPadSpec) {
       size_t N = Spec.find(':');
       if (N == std::string::npos)
         continue;
@@ -324,32 +319,6 @@ bool BinaryEmitter::emitFunction(BinaryFunction &Function,
     Streamer.emitCodeAlignment(Function.getAlign(), &*BC.STI);
   }
 
-  if (size_t Padding =
-          opts::padFunction(opts::FunctionPadBeforeSpec, Function)) {
-    // Handle padFuncsBefore after the above alignment logic but before
-    // symbol addresses are decided.
-    if (!BC.HasRelocations) {
-      BC.errs() << "BOLT-ERROR: -pad-before-funcs is not supported in "
-                << "non-relocation mode\n";
-      exit(1);
-    }
-
-    // Preserve Function.getMinAlign().
-    if (!isAligned(Function.getMinAlign(), Padding)) {
-      BC.errs() << "BOLT-ERROR: user-requested " << Padding
-                << " padding bytes before function " << Function
-                << " is not a multiple of the minimum function alignment ("
-                << Function.getMinAlign().value() << ").\n";
-      exit(1);
-    }
-
-    LLVM_DEBUG(dbgs() << "BOLT-DEBUG: padding before function " << Function
-                      << " with " << Padding << " bytes\n");
-
-    // Since the padding is not executed, it can be null bytes.
-    Streamer.emitFill(Padding, 0);
-  }
-
   MCContext &Context = Streamer.getContext();
   const MCAsmInfo *MAI = Context.getAsmInfo();
 
@@ -404,7 +373,7 @@ bool BinaryEmitter::emitFunction(BinaryFunction &Function,
   emitFunctionBody(Function, FF, /*EmitCodeOnly=*/false);
 
   // Emit padding if requested.
-  if (size_t Padding = opts::padFunction(opts::FunctionPadSpec, Function)) {
+  if (size_t Padding = opts::padFunction(Function)) {
     LLVM_DEBUG(dbgs() << "BOLT-DEBUG: padding function " << Function << " with "
                       << Padding << " bytes\n");
     Streamer.emitFill(Padding, MAI->getTextAlignFillValue());

bolt/lib/Passes/ReorderFunctions.cpp

@@ -28,9 +28,7 @@ extern cl::OptionCategory BoltOptCategory;
 extern cl::opt<unsigned> Verbosity;
 extern cl::opt<uint32_t> RandomSeed;
 
-extern size_t padFunction(const cl::list<std::string> &Spec,
-                          const bolt::BinaryFunction &Function);
-extern cl::list<std::string> FunctionPadSpec, FunctionPadBeforeSpec;
+extern size_t padFunction(const bolt::BinaryFunction &Function);
 
 extern cl::opt<bolt::ReorderFunctions::ReorderType> ReorderFunctions;
 cl::opt<bolt::ReorderFunctions::ReorderType> ReorderFunctions(
@@ -306,12 +304,8 @@ Error ReorderFunctions::runOnFunctions(BinaryContext &BC) {
                           return false;
                         if (B->isIgnored())
                           return true;
-                        const size_t PadA =
-                            opts::padFunction(opts::FunctionPadSpec, *A) +
-                            opts::padFunction(opts::FunctionPadBeforeSpec, *A);
-                        const size_t PadB =
-                            opts::padFunction(opts::FunctionPadSpec, *B) +
-                            opts::padFunction(opts::FunctionPadBeforeSpec, *B);
+                        const size_t PadA = opts::padFunction(*A);
+                        const size_t PadB = opts::padFunction(*B);
                         if (!PadA || !PadB) {
                           if (PadA)
                             return true;

clang/bindings/python/clang/cindex.py

@@ -2125,12 +2125,26 @@ def get_field_offsetof(self):
 
     def is_anonymous(self):
         """
-        Check if the record is anonymous.
+        Check whether this is a record type without a name, or a field where
+        the type is a record type without a name.
+
+        Use is_anonymous_record_decl to check whether a record is an
+        "anonymous union" as defined in the C/C++ standard.
         """
         if self.kind == CursorKind.FIELD_DECL:
             return self.type.get_declaration().is_anonymous()
         return conf.lib.clang_Cursor_isAnonymous(self)  # type: ignore [no-any-return]
 
+    def is_anonymous_record_decl(self):
+        """
+        Check if the record is an anonymous union as defined in the C/C++ standard
+        (or an "anonymous struct", the corresponding non-standard extension for
+        structs).
+        """
+        if self.kind == CursorKind.FIELD_DECL:
+            return self.type.get_declaration().is_anonymous_record_decl()
+        return conf.lib.clang_Cursor_isAnonymousRecordDecl(self)  # type: ignore [no-any-return]
+
     def is_bitfield(self):
         """
         Check if the field is a bitfield.
@@ -3902,12 +3916,13 @@ def write_main_file_to_stdout(self):
     ("clang_Cursor_getTemplateArgumentType", [Cursor, c_uint], Type),
     ("clang_Cursor_getTemplateArgumentValue", [Cursor, c_uint], c_longlong),
     ("clang_Cursor_getTemplateArgumentUnsignedValue", [Cursor, c_uint], c_ulonglong),
-    ("clang_Cursor_isAnonymous", [Cursor], bool),
-    ("clang_Cursor_isBitField", [Cursor], bool),
     ("clang_Cursor_getBinaryOpcode", [Cursor], c_int),
     ("clang_Cursor_getBriefCommentText", [Cursor], _CXString),
     ("clang_Cursor_getRawCommentText", [Cursor], _CXString),
     ("clang_Cursor_getOffsetOfField", [Cursor], c_longlong),
+    ("clang_Cursor_isAnonymous", [Cursor], bool),
+    ("clang_Cursor_isAnonymousRecordDecl", [Cursor], bool),
+    ("clang_Cursor_isBitField", [Cursor], bool),
     ("clang_Location_isInSystemHeader", [SourceLocation], bool),
     ("clang_Type_getAlignOf", [Type], c_longlong),
     ("clang_Type_getClassType", [Type], Type),

clang/bindings/python/tests/cindex/test_type.py

@@ -463,8 +463,11 @@ def test_offset(self):
             self.assertNotEqual(children[0].spelling, "typeanon")
             self.assertEqual(children[1].spelling, "typeanon")
             self.assertEqual(fields[0].kind, CursorKind.FIELD_DECL)
+            self.assertTrue(fields[0].is_anonymous())
+            self.assertFalse(fields[0].is_anonymous_record_decl())
             self.assertEqual(fields[1].kind, CursorKind.FIELD_DECL)
             self.assertTrue(fields[1].is_anonymous())
+            self.assertTrue(fields[1].is_anonymous_record_decl())
             self.assertEqual(teststruct.type.get_offset("typeanon"), f1)
             self.assertEqual(teststruct.type.get_offset("bariton"), bariton)
             self.assertEqual(teststruct.type.get_offset("foo"), foo)

clang/docs/ReleaseNotes.rst

@@ -1309,6 +1309,8 @@ Sanitizers
 Python Binding Changes
 ----------------------
 - Fixed an issue that led to crashes when calling ``Type.get_exception_specification_kind``.
+- Added binding for ``clang_Cursor_isAnonymousRecordDecl``, which allows checking if
+  a declaration is an anonymous union or anonymous struct.
 
 OpenMP Support
 --------------

clang/include/clang/Basic/BuiltinsSPIRV.td

@@ -8,7 +8,7 @@
 
 include "clang/Basic/BuiltinsBase.td"
 
-def HLSLDistance : Builtin {
+def SPIRVDistance : Builtin {
   let Spellings = ["__builtin_spirv_distance"];
   let Attributes = [NoThrow, Const];
   let Prototype = "void(...)";

clang/lib/CodeGen/CodeGenFunction.cpp

@@ -2102,7 +2102,10 @@ void CodeGenFunction::EmitBranchOnBoolExpr(
          MDHelper.createConstant(BranchHintConstant)});
     BrInst->setMetadata("hlsl.controlflow.hint",
                         llvm::MDNode::get(CGM.getLLVMContext(), Vals));
-  } break;
+    break;
+  }
+  case HLSLControlFlowHintAttr::SpellingNotCalculated:
+    break;
   }
 }
 

clang/lib/Driver/ToolChains/SYCL.cpp

@@ -17,8 +17,7 @@ using namespace llvm::opt;
 
 SYCLInstallationDetector::SYCLInstallationDetector(
     const Driver &D, const llvm::Triple &HostTriple,
-    const llvm::opt::ArgList &Args)
-    : D(D) {}
+    const llvm::opt::ArgList &Args) {}
 
 void SYCLInstallationDetector::addSYCLIncludeArgs(
     const ArgList &DriverArgs, ArgStringList &CC1Args) const {
@@ -31,8 +30,8 @@ void SYCLInstallationDetector::addSYCLIncludeArgs(
 }
 
 // Unsupported options for SYCL device compilation.
-static ArrayRef<OptSpecifier> getUnsupportedOpts() {
-  return {
+static ArrayRef<options::ID> getUnsupportedOpts() {
+  static constexpr options::ID UnsupportedOpts[] = {
       options::OPT_fsanitize_EQ,      // -fsanitize
       options::OPT_fcf_protection_EQ, // -fcf-protection
       options::OPT_fprofile_generate,
@@ -53,7 +52,9 @@ static ArrayRef<OptSpecifier> getUnsupportedOpts() {
       options::OPT_fprofile_instr_use_EQ,       // -fprofile-instr-use
       options::OPT_forder_file_instrumentation, // -forder-file-instrumentation
       options::OPT_fcs_profile_generate,        // -fcs-profile-generate
-      options::OPT_fcs_profile_generate_EQ};
+      options::OPT_fcs_profile_generate_EQ,
+  };
+  return UnsupportedOpts;
 }
 
 SYCLToolChain::SYCLToolChain(const Driver &D, const llvm::Triple &Triple,

clang/lib/Driver/ToolChains/SYCL.h

@@ -22,9 +22,6 @@ class SYCLInstallationDetector {
 
   void addSYCLIncludeArgs(const llvm::opt::ArgList &DriverArgs,
                           llvm::opt::ArgStringList &CC1Args) const;
-
-private:
-  const Driver &D;
 };
 
 namespace toolchains {

clang/test/Driver/print-supported-extensions-riscv.c

@@ -193,6 +193,7 @@
 // CHECK-NEXT:     xqcia                0.2       'Xqcia' (Qualcomm uC Arithmetic Extension)
 // CHECK-NEXT:     xqciac               0.2       'Xqciac' (Qualcomm uC Load-Store Address Calculation Extension)
 // CHECK-NEXT:     xqcicli              0.2       'Xqcicli' (Qualcomm uC Conditional Load Immediate Extension)
+// CHECK-NEXT:     xqcicm               0.2       'Xqcicm' (Qualcomm uC Conditional Move Extension)
 // CHECK-NEXT:     xqcics               0.2       'Xqcics' (Qualcomm uC Conditional Select Extension)
 // CHECK-NEXT:     xqcicsr              0.2       'Xqcicsr' (Qualcomm uC CSR Extension)
 // CHECK-NEXT:     xqcilsm              0.2       'Xqcilsm' (Qualcomm uC Load Store Multiple Extension)

clang/test/Driver/sycl-offload-jit.cpp

@@ -3,7 +3,7 @@
 /// Check the phases graph with -fsycl. Use of -fsycl enables offload
 // RUN: %clang -ccc-print-phases --target=x86_64-unknown-linux-gnu -fsycl %s 2>&1 \
 // RUN:   | FileCheck -check-prefixes=CHK-PHASES %s
-// RUN: %clang_cl -ccc-print-phases --target=x86_64-pc-windows-msvc -fsycl %s 2>&1 \
+// RUN: %clang_cl -ccc-print-phases --target=x86_64-pc-windows-msvc -fsycl -- %s 2>&1 \
 // RUN:   | FileCheck -check-prefixes=CHK-PHASES %s
 // CHK-PHASES: 0: input, "[[INPUT:.+\.cpp]]", c++, (host-sycl)
 // CHK-PHASES-NEXT: 1: preprocessor, {0}, c++-cpp-output, (host-sycl)
@@ -35,7 +35,7 @@
 // RUN:   | FileCheck -check-prefixes=CHK-FSYCL-IS-DEVICE,CHK-FSYCL-IS-HOST %s
 // RUN: %clang -### -fsycl -fsycl-device-only %s 2>&1 \
 // RUN:   | FileCheck -check-prefix=CHK-FSYCL-IS-DEVICE %s
-// RUN: %clang_cl -### -fsycl -c %s 2>&1 \
+// RUN: %clang_cl -### -fsycl -c -- %s 2>&1 \
 // RUN:   | FileCheck -check-prefixes=CHK-FSYCL-IS-DEVICE,CHK-FSYCL-IS-HOST %s
 // RUN: %clang -### -fsycl -fsycl-host-only %s 2>&1 \
 // RUN:   | FileCheck -check-prefix=CHK-FSYCL-IS-HOST %s

flang/lib/Common/Fortran.cpp

@@ -136,7 +136,8 @@ bool AreCompatibleCUDADataAttrs(std::optional<CUDADataAttr> x,
       if (*x == CUDADataAttr::Device) {
         if ((y &&
                 (*y == CUDADataAttr::Managed || *y == CUDADataAttr::Unified ||
-                    *y == CUDADataAttr::Shared)) ||
+                    *y == CUDADataAttr::Shared ||
+                    *y == CUDADataAttr::Constant)) ||
             (!y && (isCudaUnified || isCudaManaged))) {
           if (y && *y == CUDADataAttr::Shared && warning) {
             *warning = "SHARED attribute ignored"s;

flang/test/Driver/parse-error.ll

@@ -13,7 +13,7 @@
 ; RUN: not %flang_fc1 -fdebug-unparse-no-sema -x f95 %s 2>&1 | FileCheck %s --check-prefix=ERROR
 ; RUN: not %flang_fc1 -fsyntax-only %s -x f95 2>&1 | FileCheck %s --check-prefix=ERROR
 
-; ERROR: Could not scan {{.*}}parse-error.f95
+; ERROR: Could not scan {{.*}}parse-error.ll
 
 define void @foo() {
   ret void

flang/test/Semantics/cuf10.cuf

@@ -2,6 +2,7 @@
 module m
   real, device :: a(4,8)
   real, managed, allocatable :: b(:,:)
+  integer, constant :: x = 1
  contains
   attributes(global) subroutine kernel(a,b,c,n,m)
     integer, value :: n
@@ -23,4 +24,10 @@ module m
       call devsub(c,4) ! not checked in OpenACC construct
     end do
   end
+  attributes(global) subroutine sub1(x)
+    integer :: x
+  end
+  subroutine sub2()
+    call sub1<<<1,1>>>(x) ! actual constant to device dummy
+  end
 end