diff --git a/clang-tools-extra/clang-tidy/modernize/UseDesignatedInitializersCheck.cpp b/clang-tools-extra/clang-tidy/modernize/UseDesignatedInitializersCheck.cpp
index 2a0cc403b726e8..3132067f3d5ece 100644
--- a/clang-tools-extra/clang-tidy/modernize/UseDesignatedInitializersCheck.cpp
+++ b/clang-tools-extra/clang-tidy/modernize/UseDesignatedInitializersCheck.cpp
@@ -80,9 +80,13 @@ unsigned getNumberOfDesignated(const InitListExpr *SyntacticInitList) {
   });
 }
 
-AST_MATCHER(CXXRecordDecl, isAggregate) { return Node.isAggregate(); }
+AST_MATCHER(CXXRecordDecl, isAggregate) {
+  return Node.hasDefinition() && Node.isAggregate();
+}
 
-AST_MATCHER(CXXRecordDecl, isPOD) { return Node.isPOD(); }
+AST_MATCHER(CXXRecordDecl, isPOD) {
+  return Node.hasDefinition() && Node.isPOD();
+}
 
 AST_MATCHER(InitListExpr, isFullyDesignated) {
   if (const InitListExpr *SyntacticForm =
diff --git a/clang-tools-extra/docs/ReleaseNotes.rst b/clang-tools-extra/docs/ReleaseNotes.rst
index 54118e5f92f417..ccebf74e8a67e7 100644
--- a/clang-tools-extra/docs/ReleaseNotes.rst
+++ b/clang-tools-extra/docs/ReleaseNotes.rst
@@ -216,6 +216,10 @@ Changes in existing checks
   a false positive when only an implicit conversion happened inside an
   initializer list.
 
+- Improved :doc:`modernize-use-designated-initializers
+  <clang-tidy/checks/modernize/use-designated-initializers>` check to fix a
+  crash when a class is declared but not defined.
+
 - Improved :doc:`modernize-use-nullptr
   <clang-tidy/checks/modernize/use-nullptr>` check to also recognize
   ``NULL``/``__null`` (but not ``0``) when used with a templated type.
diff --git a/clang-tools-extra/test/clang-tidy/checkers/modernize/use-designated-initializers.cpp b/clang-tools-extra/test/clang-tidy/checkers/modernize/use-designated-initializers.cpp
index 9b769ad0be23ca..048665b2e54ac5 100644
--- a/clang-tools-extra/test/clang-tidy/checkers/modernize/use-designated-initializers.cpp
+++ b/clang-tools-extra/test/clang-tidy/checkers/modernize/use-designated-initializers.cpp
@@ -201,3 +201,11 @@ DECLARE_S93;
 // CHECK-MESSAGES-MACROS: :[[@LINE-1]]:1: warning: use designated initializer list to initialize 'S9' [modernize-use-designated-initializers]
 // CHECK-MESSAGES-MACROS: :[[@LINE-4]]:28: note: expanded from macro 'DECLARE_S93'
 // CHECK-MESSAGES-MACROS: :[[@LINE-71]]:1: note: aggregate type is defined here
+
+// Issue #113652.
+struct S14;
+
+struct S15{
+  S15(S14& d):d{d}{}
+  S14& d;
+};
diff --git a/compiler-rt/test/asan/TestCases/Posix/ignore_free_hook.cpp b/compiler-rt/test/asan/TestCases/Posix/ignore_free_hook.cpp
index 87be90014d56e8..dfeb8ad5c7b53f 100644
--- a/compiler-rt/test/asan/TestCases/Posix/ignore_free_hook.cpp
+++ b/compiler-rt/test/asan/TestCases/Posix/ignore_free_hook.cpp
@@ -26,14 +26,17 @@ bool ignore_free = false;
 
 extern "C" {
 WEAK_ON_APPLE void __sanitizer_free_hook(const volatile void *ptr) {
-  if (ptr == glob_ptr)
+  if (ptr == glob_ptr) {
     fprintf(stderr, "Free Hook\n");
+    fflush(stderr);
+  }
 }
 
 WEAK_ON_APPLE int __sanitizer_ignore_free_hook(const volatile void *ptr) {
   if (ptr != glob_ptr)
     return 0;
   fprintf(stderr, ignore_free ? "Free Ignored\n" : "Free Respected\n");
+  fflush(stderr);
   return ignore_free;
 }
 } // extern "C"
diff --git a/compiler-rt/test/hwasan/TestCases/many-threads-uaf.c b/compiler-rt/test/hwasan/TestCases/many-threads-uaf.c
index 8fa07861371d56..e02ab5b28ce046 100644
--- a/compiler-rt/test/hwasan/TestCases/many-threads-uaf.c
+++ b/compiler-rt/test/hwasan/TestCases/many-threads-uaf.c
@@ -23,6 +23,7 @@ void *BoringThread(void *arg) {
 void *UAFThread(void *arg) {
   char * volatile x = (char*)malloc(10);
   fprintf(stderr, "ZZZ %p\n", x);
+  fflush(stderr);
   free(x);
   x[5] = 42;
   // CHECK: ERROR: HWAddressSanitizer: tag-mismatch on address
diff --git a/compiler-rt/test/hwasan/TestCases/mem-intrinsics.c b/compiler-rt/test/hwasan/TestCases/mem-intrinsics.c
index 78bef538af1161..da1cb686969206 100644
--- a/compiler-rt/test/hwasan/TestCases/mem-intrinsics.c
+++ b/compiler-rt/test/hwasan/TestCases/mem-intrinsics.c
@@ -21,6 +21,7 @@ int main() {
   memcpy(Q, P, 32);
 #endif
   write(STDOUT_FILENO, "recovered\n", 10);
+  fflush(stdout);
   // WRITE: ERROR: HWAddressSanitizer: tag-mismatch on address
   // WRITE: WRITE of size 32 at {{.*}} tags: [[PTR_TAG:..]]/[[MEM_TAG:..]] (ptr/mem)
   // WRITE: Invalid access starting at offset 16
diff --git a/compiler-rt/test/hwasan/TestCases/use-after-free.c b/compiler-rt/test/hwasan/TestCases/use-after-free.c
index 070622f560a225..b4b79875e8111e 100644
--- a/compiler-rt/test/hwasan/TestCases/use-after-free.c
+++ b/compiler-rt/test/hwasan/TestCases/use-after-free.c
@@ -15,6 +15,7 @@ int main() {
   free(x);
   __hwasan_disable_allocator_tagging();
   fprintf(stderr, ISREAD ? "Going to do a READ\n" : "Going to do a WRITE\n");
+  fflush(stderr);
   // CHECK: Going to do a [[TYPE:[A-Z]*]]
   int r = 0;
   if (ISREAD) r = x[5]; else x[5] = 42;  // should be on the same line.
diff --git a/flang/lib/Optimizer/Transforms/CUFAddConstructor.cpp b/flang/lib/Optimizer/Transforms/CUFAddConstructor.cpp
index 4da06be8ef7dd9..7cdb2f7ffe27d9 100644
--- a/flang/lib/Optimizer/Transforms/CUFAddConstructor.cpp
+++ b/flang/lib/Optimizer/Transforms/CUFAddConstructor.cpp
@@ -6,15 +6,23 @@
 //
 //===----------------------------------------------------------------------===//
 
+#include "flang/Optimizer/Builder/BoxValue.h"
 #include "flang/Optimizer/Builder/FIRBuilder.h"
+#include "flang/Optimizer/Builder/Runtime/RTBuilder.h"
+#include "flang/Optimizer/Builder/Todo.h"
+#include "flang/Optimizer/CodeGen/Target.h"
 #include "flang/Optimizer/Dialect/CUF/CUFOps.h"
 #include "flang/Optimizer/Dialect/FIRAttr.h"
 #include "flang/Optimizer/Dialect/FIRDialect.h"
+#include "flang/Optimizer/Dialect/FIROps.h"
 #include "flang/Optimizer/Dialect/FIROpsSupport.h"
+#include "flang/Optimizer/Support/DataLayout.h"
 #include "flang/Optimizer/Transforms/CUFCommon.h"
+#include "flang/Runtime/CUDA/registration.h"
 #include "flang/Runtime/entry-names.h"
 #include "mlir/Dialect/GPU/IR/GPUDialect.h"
 #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
+#include "mlir/IR/Value.h"
 #include "mlir/Pass/Pass.h"
 #include "llvm/ADT/SmallVector.h"
 
@@ -23,6 +31,8 @@ namespace fir {
 #include "flang/Optimizer/Transforms/Passes.h.inc"
 } // namespace fir
 
+using namespace Fortran::runtime::cuda;
+
 namespace {
 
 static constexpr llvm::StringRef cudaFortranCtorName{
@@ -34,13 +44,23 @@ struct CUFAddConstructor
   void runOnOperation() override {
     mlir::ModuleOp mod = getOperation();
     mlir::SymbolTable symTab(mod);
-    mlir::OpBuilder builder{mod.getBodyRegion()};
+    mlir::OpBuilder opBuilder{mod.getBodyRegion()};
+    fir::FirOpBuilder builder(opBuilder, mod);
+    fir::KindMapping kindMap{fir::getKindMapping(mod)};
     builder.setInsertionPointToEnd(mod.getBody());
     mlir::Location loc = mod.getLoc();
     auto *ctx = mod.getContext();
     auto voidTy = mlir::LLVM::LLVMVoidType::get(ctx);
+    auto idxTy = builder.getIndexType();
     auto funcTy =
         mlir::LLVM::LLVMFunctionType::get(voidTy, {}, /*isVarArg=*/false);
+    std::optional<mlir::DataLayout> dl =
+        fir::support::getOrSetDataLayout(mod, /*allowDefaultLayout=*/false);
+    if (!dl) {
+      mlir::emitError(mod.getLoc(),
+                      "data layout attribute is required to perform " +
+                          getName() + "pass");
+    }
 
     // Symbol reference to CUFRegisterAllocator.
     builder.setInsertionPointToEnd(mod.getBody());
@@ -58,12 +78,13 @@ struct CUFAddConstructor
     builder.setInsertionPointToStart(func.addEntryBlock(builder));
     builder.create<mlir::LLVM::CallOp>(loc, funcTy, cufRegisterAllocatorRef);
 
-    // Register kernels
     auto gpuMod = symTab.lookup<mlir::gpu::GPUModuleOp>(cudaDeviceModuleName);
     if (gpuMod) {
       auto llvmPtrTy = mlir::LLVM::LLVMPointerType::get(ctx);
       auto registeredMod = builder.create<cuf::RegisterModuleOp>(
           loc, llvmPtrTy, mlir::SymbolRefAttr::get(ctx, gpuMod.getName()));
+
+      // Register kernels
       for (auto func : gpuMod.getOps<mlir::gpu::GPUFuncOp>()) {
         if (func.isKernel()) {
           auto kernelName = mlir::SymbolRefAttr::get(
@@ -72,12 +93,55 @@ struct CUFAddConstructor
           builder.create<cuf::RegisterKernelOp>(loc, kernelName, registeredMod);
         }
       }
+
+      // Register variables
+      for (fir::GlobalOp globalOp : mod.getOps<fir::GlobalOp>()) {
+        auto attr = globalOp.getDataAttrAttr();
+        if (!attr)
+          continue;
+
+        mlir::func::FuncOp func;
+        switch (attr.getValue()) {
+        case cuf::DataAttribute::Device:
+        case cuf::DataAttribute::Constant: {
+          func = fir::runtime::getRuntimeFunc<mkRTKey(CUFRegisterVariable)>(
+              loc, builder);
+          auto fTy = func.getFunctionType();
+
+          // Global variable name
+          std::string gblNameStr = globalOp.getSymbol().getValue().str();
+          gblNameStr += '\0';
+          mlir::Value gblName = fir::getBase(
+              fir::factory::createStringLiteral(builder, loc, gblNameStr));
+
+          // Global variable size
+          auto sizeAndAlign = fir::getTypeSizeAndAlignmentOrCrash(
+              loc, globalOp.getType(), *dl, kindMap);
+          auto size =
+              builder.createIntegerConstant(loc, idxTy, sizeAndAlign.first);
+
+          // Global variable address
+          mlir::Value addr = builder.create<fir::AddrOfOp>(
+              loc, globalOp.resultType(), globalOp.getSymbol());
+
+          llvm::SmallVector<mlir::Value> args{fir::runtime::createArguments(
+              builder, loc, fTy, registeredMod, addr, gblName, size)};
+          builder.create<fir::CallOp>(loc, func, args);
+        } break;
+        case cuf::DataAttribute::Managed:
+          TODO(loc, "registration of managed variables");
+        default:
+          break;
+        }
+        if (!func)
+          continue;
+      }
     }
     builder.create<mlir::LLVM::ReturnOp>(loc, mlir::ValueRange{});
 
     // Create the llvm.global_ctor with the function.
-    // TODO: We might want to have a utility that retrieve it if already created
-    // and adds new functions.
+    // TODO: We might want to have a utility that retrieve it if already
+    // created and adds new functions.
     builder.setInsertionPointToEnd(mod.getBody());
     llvm::SmallVector<mlir::Attribute> funcs;
     funcs.push_back(
diff --git a/flang/lib/Optimizer/Transforms/CUFOpConversion.cpp b/flang/lib/Optimizer/Transforms/CUFOpConversion.cpp
index 9c2b882c7f46fe..14cc1cb508cfc0 100644
--- a/flang/lib/Optimizer/Transforms/CUFOpConversion.cpp
+++ b/flang/lib/Optimizer/Transforms/CUFOpConversion.cpp
@@ -111,7 +111,7 @@ mlir::Value getDeviceAddress(mlir::PatternRewriter &rewriter,
     switch (attr.getValue()) {
     case cuf::DataAttribute::Device:
     case cuf::DataAttribute::Managed:
-    case cuf::DataAttribute::Pinned:
+    case cuf::DataAttribute::Constant:
       isDevGlobal = true;
       break;
     default:
diff --git a/flang/test/Fir/CUDA/cuda-constructor-2.f90 b/flang/test/Fir/CUDA/cuda-constructor-2.f90
new file mode 100644
index 00000000000000..378dabbb7c7e7d
--- /dev/null
+++ b/flang/test/Fir/CUDA/cuda-constructor-2.f90
@@ -0,0 +1,22 @@
+// RUN: fir-opt --split-input-file --cuf-add-constructor %s | FileCheck %s
+
+module attributes {dlti.dl_spec = #dlti.dl_spec<#dlti.dl_entry<!llvm.ptr, dense<64> : vector<4xi64>>, #dlti.dl_entry<!llvm.ptr<271>, dense<32> : vector<4xi64>>, #dlti.dl_entry<!llvm.ptr<270>, dense<32> : vector<4xi64>>, #dlti.dl_entry<f128, dense<128> : vector<2xi64>>, #dlti.dl_entry<f64, dense<64> : vector<2xi64>>, #dlti.dl_entry<f80, dense<128> : vector<2xi64>>, #dlti.dl_entry<f16, dense<16> : vector<2xi64>>, #dlti.dl_entry<i32, dense<32> : vector<2xi64>>, #dlti.dl_entry<i16, dense<16> : vector<2xi64>>, #dlti.dl_entry<i128, dense<128> : vector<2xi64>>, #dlti.dl_entry<i8, dense<8> : vector<2xi64>>, #dlti.dl_entry<!llvm.ptr<272>, dense<64> : vector<4xi64>>, #dlti.dl_entry<i64, dense<64> : vector<2xi64>>, #dlti.dl_entry<i1, dense<8> : vector<2xi64>>, #dlti.dl_entry<"dlti.endianness", "little">, #dlti.dl_entry<"dlti.stack_alignment", 128 : i64>>, fir.defaultkind = "a1c4d8i4l4r4", fir.kindmap = "", gpu.container_module, llvm.data_layout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-i128:128-f80:128-n8:16:32:64-S128", llvm.ident = "flang version 20.0.0 (https://github.com/llvm/llvm-project.git cae351f3453a0a26ec8eb2ddaf773c24a29d929e)", llvm.target_triple = "x86_64-unknown-linux-gnu"} {
+
+  fir.global @_QMmtestsEn(dense<[3, 4, 5, 6, 7]> : tensor<5xi32>) {data_attr = #cuf.cuda<device>} : !fir.array<5xi32>
+
+  gpu.module @cuda_device_mod [#nvvm.target] {
+  }
+}
+
+// CHECK: gpu.module @cuda_device_mod [#nvvm.target] 
+
+// CHECK: llvm.func internal @__cudaFortranConstructor() {
+// CHECK-DAG: %[[MODULE:.*]] = cuf.register_module @cuda_device_mod -> !llvm.ptr
+// CHECK-DAG: %[[VAR_NAME:.*]] = fir.address_of(@_QQ{{.*}}) : !fir.ref<!fir.char<1,12>>
+// CHECK-DAG: %[[VAR_ADDR:.*]] = fir.address_of(@_QMmtestsEn) : !fir.ref<!fir.array<5xi32>>
+// CHECK-DAG: %[[MODULE2:.*]] = fir.convert %[[MODULE]] : (!llvm.ptr) -> !fir.ref<!fir.llvm_ptr<i8>>
+// CHECK-DAG: %[[VAR_ADDR2:.*]] = fir.convert %[[VAR_ADDR]] : (!fir.ref<!fir.array<5xi32>>) -> !fir.ref<i8>
+// CHECK-DAG: %[[VAR_NAME2:.*]] = fir.convert %[[VAR_NAME]] : (!fir.ref<!fir.char<1,12>>) -> !fir.ref<i8>
+// CHECK-DAG: %[[CST:.*]] = arith.constant 20 : index
+// CHECK-DAG %[[CST2:.*]] = fir.convert %[[CST]] : (index) -> i64
+// CHECK fir.call @_FortranACUFRegisterVariable(%[[MODULE2]], %[[VAR_ADDR2]], %[[VAR_NAME2]], %[[CST2]]) : (!fir.ref<!fir.llvm_ptr<i8>>, !fir.ref<i8>, !fir.ref<i8>, i64) -> none
diff --git a/llvm/include/llvm/CodeGen/GlobalISel/GIMatchTableExecutor.h b/llvm/include/llvm/CodeGen/GlobalISel/GIMatchTableExecutor.h
index 7b42722ca8d4f1..b4ff4cd178d757 100644
--- a/llvm/include/llvm/CodeGen/GlobalISel/GIMatchTableExecutor.h
+++ b/llvm/include/llvm/CodeGen/GlobalISel/GIMatchTableExecutor.h
@@ -24,6 +24,7 @@
 #include "llvm/CodeGen/MachineInstr.h"
 #include "llvm/CodeGenTypes/LowLevelType.h"
 #include "llvm/IR/Function.h"
+#include "llvm/Transforms/Utils/SizeOpts.h"
 #include <bitset>
 #include <cstddef>
 #include <cstdint>
@@ -635,8 +636,12 @@ class GIMatchTableExecutor {
 
   bool shouldOptForSize(const MachineFunction *MF) const {
     const auto &F = MF->getFunction();
-    return F.hasOptSize() || F.hasMinSize() ||
-           (PSI && BFI && CurMBB && llvm::shouldOptForSize(*CurMBB, PSI, BFI));
+    if (F.hasOptSize())
+      return true;
+    if (CurMBB)
+      if (auto *BB = CurMBB->getBasicBlock())
+        return llvm::shouldOptimizeForSize(BB, PSI, BFI);
+    return false;
   }
 
 public:
diff --git a/llvm/include/llvm/CodeGen/GlobalISel/Utils.h b/llvm/include/llvm/CodeGen/GlobalISel/Utils.h
index 95a8234d3c6080..4016247376c4f6 100644
--- a/llvm/include/llvm/CodeGen/GlobalISel/Utils.h
+++ b/llvm/include/llvm/CodeGen/GlobalISel/Utils.h
@@ -542,10 +542,6 @@ bool isConstFalseVal(const TargetLowering &TLI, int64_t Val, bool IsVector,
 /// TargetBooleanContents.
 int64_t getICmpTrueVal(const TargetLowering &TLI, bool IsVector, bool IsFP);
 
-/// Returns true if the given block should be optimized for size.
-bool shouldOptForSize(const MachineBasicBlock &MBB, ProfileSummaryInfo *PSI,
-                      BlockFrequencyInfo *BFI);
-
 using SmallInstListTy = GISelWorkList<4>;
 void saveUsesAndErase(MachineInstr &MI, MachineRegisterInfo &MRI,
                       LostDebugLocObserver *LocObserver,
diff --git a/llvm/include/llvm/IR/IntrinsicsSPIRV.td b/llvm/include/llvm/IR/IntrinsicsSPIRV.td
index 6df2eb156a0774..ddb47390537412 100644
--- a/llvm/include/llvm/IR/IntrinsicsSPIRV.td
+++ b/llvm/include/llvm/IR/IntrinsicsSPIRV.td
@@ -87,6 +87,7 @@ let TargetPrefix = "spv" in {
   def int_spv_wave_readlane : DefaultAttrsIntrinsic<[llvm_any_ty], [LLVMMatchType<0>, llvm_i32_ty], [IntrConvergent, IntrNoMem]>;
   def int_spv_sign : DefaultAttrsIntrinsic<[LLVMScalarOrSameVectorWidth<0, llvm_i32_ty>], [llvm_any_ty], [IntrNoMem]>;
   def int_spv_radians : DefaultAttrsIntrinsic<[LLVMMatchType<0>], [llvm_anyfloat_ty], [IntrNoMem]>;
+  def int_spv_group_memory_barrier_with_group_sync : DefaultAttrsIntrinsic<[], [], []>;
 
   // Create resource handle given the binding information. Returns a 
   // type appropriate for the kind of resource given the set id, binding id,
diff --git a/llvm/lib/AsmParser/LLLexer.cpp b/llvm/lib/AsmParser/LLLexer.cpp
index 759db6db60774c..56abd03d623541 100644
--- a/llvm/lib/AsmParser/LLLexer.cpp
+++ b/llvm/lib/AsmParser/LLLexer.cpp
@@ -60,8 +60,8 @@ uint64_t LLLexer::atoull(const char *Buffer, const char *End) {
     uint64_t OldRes = Result;
     Result *= 10;
     Result += *Buffer-'0';
-    if (Result < OldRes) {  // Uh, oh, overflow detected!!!
-      LexError("constant bigger than 64 bits detected!");
+    if (Result < OldRes) { // overflow detected.
+      LexError("constant bigger than 64 bits detected");
       return 0;
     }
   }
@@ -75,8 +75,8 @@ uint64_t LLLexer::HexIntToVal(const char *Buffer, const char *End) {
     Result *= 16;
     Result += hexDigitValue(*Buffer);
 
-    if (Result < OldRes) {   // Uh, oh, overflow detected!!!
-      LexError("constant bigger than 64 bits detected!");
+    if (Result < OldRes) { // overflow detected.
+      LexError("constant bigger than 64 bits detected");
       return 0;
     }
   }
@@ -99,7 +99,7 @@ void LLLexer::HexToIntPair(const char *Buffer, const char *End,
     Pair[1] += hexDigitValue(*Buffer);
   }
   if (Buffer != End)
-    LexError("constant bigger than 128 bits detected!");
+    LexError("constant bigger than 128 bits detected");
 }
 
 /// FP80HexToIntPair - translate an 80 bit FP80 number (20 hexits) into
@@ -118,7 +118,7 @@ void LLLexer::FP80HexToIntPair(const char *Buffer, const char *End,
     Pair[0] += hexDigitValue(*Buffer);
   }
   if (Buffer != End)
-    LexError("constant bigger than 128 bits detected!");
+    LexError("constant bigger than 128 bits detected");
 }
 
 // UnEscapeLexed - Run through the specified buffer and change \xx codes to the
@@ -292,7 +292,7 @@ lltok::Kind LLLexer::LexDollar() {
         StrVal.assign(TokStart + 2, CurPtr - 1);
         UnEscapeLexed(StrVal);
         if (StringRef(StrVal).contains(0)) {
-          LexError("Null bytes are not allowed in names");
+          LexError("NUL character is not allowed in names");
           return lltok::Error;
         }
         return lltok::ComdatVar;
@@ -354,7 +354,7 @@ lltok::Kind LLLexer::LexUIntID(lltok::Kind Token) {
 
   uint64_t Val = atoull(TokStart + 1, CurPtr);
   if ((unsigned)Val != Val)
-    LexError("invalid value number (too large)!");
+    LexError("invalid value number (too large)");
   UIntVal = unsigned(Val);
   return Token;
 }
@@ -375,7 +375,7 @@ lltok::Kind LLLexer::LexVar(lltok::Kind Var, lltok::Kind VarID) {
         StrVal.assign(TokStart+2, CurPtr-1);
         UnEscapeLexed(StrVal);
         if (StringRef(StrVal).contains(0)) {
-          LexError("Null bytes are not allowed in names");
+          LexError("NUL character is not allowed in names");
           return lltok::Error;
         }
         return Var;
@@ -410,7 +410,7 @@ lltok::Kind LLLexer::LexQuote() {
   if (CurPtr[0] == ':') {
     ++CurPtr;
     if (StringRef(StrVal).contains(0)) {
-      LexError("Null bytes are not allowed in names");
+      LexError("NUL character is not allowed in names");
       kind = lltok::Error;
     } else {
       kind = lltok::LabelStr;
@@ -492,7 +492,7 @@ lltok::Kind LLLexer::LexIdentifier() {
     uint64_t NumBits = atoull(StartChar, CurPtr);
     if (NumBits < IntegerType::MIN_INT_BITS ||
         NumBits > IntegerType::MAX_INT_BITS) {
-      LexError("bitwidth for integer type out of range!");
+      LexError("bitwidth for integer type out of range");
       return lltok::Error;
     }
     TyVal = IntegerType::get(Context, NumBits);
@@ -1122,7 +1122,7 @@ lltok::Kind LLLexer::LexDigitOrNegative() {
     uint64_t Val = atoull(TokStart, CurPtr);
     ++CurPtr; // Skip the colon.
     if ((unsigned)Val != Val)
-      LexError("invalid value number (too large)!");
+      LexError("invalid value number (too large)");
     UIntVal = unsigned(Val);
     return lltok::LabelID;
   }
diff --git a/llvm/lib/CodeGen/GlobalISel/Utils.cpp b/llvm/lib/CodeGen/GlobalISel/Utils.cpp
index 513a49b4fc2e4d..dcbbb0871a8445 100644
--- a/llvm/lib/CodeGen/GlobalISel/Utils.cpp
+++ b/llvm/lib/CodeGen/GlobalISel/Utils.cpp
@@ -1619,11 +1619,6 @@ int64_t llvm::getICmpTrueVal(const TargetLowering &TLI, bool IsVector,
   llvm_unreachable("Invalid boolean contents");
 }
 
-bool llvm::shouldOptForSize(const MachineBasicBlock &MBB,
-                            ProfileSummaryInfo *PSI, BlockFrequencyInfo *BFI) {
-  return llvm::shouldOptimizeForSize(MBB.getBasicBlock(), PSI, BFI);
-}
-
 void llvm::saveUsesAndErase(MachineInstr &MI, MachineRegisterInfo &MRI,
                             LostDebugLocObserver *LocObserver,
                             SmallInstListTy &DeadInstChain) {
diff --git a/llvm/lib/SandboxIR/Context.cpp b/llvm/lib/SandboxIR/Context.cpp
index 5e5cbbbc4515d2..b86ed5864c1ac1 100644
--- a/llvm/lib/SandboxIR/Context.cpp
+++ b/llvm/lib/SandboxIR/Context.cpp
@@ -681,7 +681,7 @@ void Context::runMoveInstrCallbacks(Instruction *I, const BBIterator &WhereIt) {
 // An arbitrary limit, to check for accidental misuse. We expect a small number
 // of callbacks to be registered at a time, but we can increase this number if
 // we discover we needed more.
-static constexpr int MaxRegisteredCallbacks = 16;
+[[maybe_unused]] static constexpr int MaxRegisteredCallbacks = 16;
 
 Context::CallbackID Context::registerEraseInstrCallback(EraseInstrCallback CB) {
   assert(EraseInstrCallbacks.size() <= MaxRegisteredCallbacks &&
diff --git a/llvm/lib/TableGen/TGLexer.cpp b/llvm/lib/TableGen/TGLexer.cpp
index 8fe7f69ecf8e59..1e93b2c160ba58 100644
--- a/llvm/lib/TableGen/TGLexer.cpp
+++ b/llvm/lib/TableGen/TGLexer.cpp
@@ -89,7 +89,7 @@ TGLexer::TGLexer(SourceMgr &SM, ArrayRef<std::string> Macros) : SrcMgr(SM) {
   for (StringRef MacroName : Macros) {
     const char *End = lexMacroName(MacroName);
     if (End != MacroName.end())
-      PrintFatalError("Invalid macro name `" + MacroName +
+      PrintFatalError("invalid macro name `" + MacroName +
                       "` specified on command line");
 
     DefinedMacros.insert(MacroName);
@@ -188,7 +188,7 @@ tgtok::TokKind TGLexer::LexToken(bool FileOrLineStart) {
       return LexIdentifier();
 
     // Unknown character, emit an error.
-    return ReturnError(TokStart, "Unexpected character");
+    return ReturnError(TokStart, "unexpected character");
   case EOF:
     // Lex next token, if we just left an include file.
     // Note that leaving an include file means that the next
@@ -231,7 +231,7 @@ tgtok::TokKind TGLexer::LexToken(bool FileOrLineStart) {
         ++CurPtr; // Eat third dot.
         return tgtok::dotdotdot;
       }
-      return ReturnError(TokStart, "Invalid '..' punctuation");
+      return ReturnError(TokStart, "invalid '..' punctuation");
     }
     return tgtok::dot;
 
@@ -255,7 +255,7 @@ tgtok::TokKind TGLexer::LexToken(bool FileOrLineStart) {
       if (SkipCComment())
         return tgtok::Error;
     } else // Otherwise, this is an error.
-      return ReturnError(TokStart, "Unexpected character");
+      return ReturnError(TokStart, "unexpected character");
     return LexToken(FileOrLineStart);
   case '-': case '+':
   case '0': case '1': case '2': case '3': case '4': case '5': case '6':
@@ -313,10 +313,10 @@ tgtok::TokKind TGLexer::LexString() {
   while (*CurPtr != '"') {
     // If we hit the end of the buffer, report an error.
     if (*CurPtr == 0 && CurPtr == CurBuf.end())
-      return ReturnError(StrStart, "End of file in string literal");
+      return ReturnError(StrStart, "end of file in string literal");
 
     if (*CurPtr == '\n' || *CurPtr == '\r')
-      return ReturnError(StrStart, "End of line in string literal");
+      return ReturnError(StrStart, "end of line in string literal");
 
     if (*CurPtr != '\\') {
       CurStrVal += *CurPtr++;
@@ -346,7 +346,7 @@ tgtok::TokKind TGLexer::LexString() {
     // If we hit the end of the buffer, report an error.
     case '\0':
       if (CurPtr == CurBuf.end())
-        return ReturnError(StrStart, "End of file in string literal");
+        return ReturnError(StrStart, "end of file in string literal");
       [[fallthrough]];
     default:
       return ReturnError(CurPtr, "invalid escape in string literal");
@@ -359,7 +359,7 @@ tgtok::TokKind TGLexer::LexString() {
 
 tgtok::TokKind TGLexer::LexVarName() {
   if (!isValidIDChar(CurPtr[0], /*First=*/true))
-    return ReturnError(TokStart, "Invalid variable name");
+    return ReturnError(TokStart, "invalid variable name");
 
   // Otherwise, we're ok, consume the rest of the characters.
   const char *VarNameStart = CurPtr++;
@@ -433,7 +433,7 @@ bool TGLexer::LexInclude() {
   tgtok::TokKind Tok = LexToken();
   if (Tok == tgtok::Error) return true;
   if (Tok != tgtok::StrVal) {
-    PrintError(getLoc(), "Expected filename after include");
+    PrintError(getLoc(), "expected filename after include");
     return true;
   }
 
@@ -444,7 +444,7 @@ bool TGLexer::LexInclude() {
   CurBuffer = SrcMgr.AddIncludeFile(Filename, SMLoc::getFromPointer(CurPtr),
                                     IncludedFile);
   if (!CurBuffer) {
-    PrintError(getLoc(), "Could not find include file '" + Filename + "'");
+    PrintError(getLoc(), "could not find include file '" + Filename + "'");
     return true;
   }
 
@@ -476,7 +476,7 @@ bool TGLexer::SkipCComment() {
     int CurChar = getNextChar();
     switch (CurChar) {
     case EOF:
-      PrintError(TokStart, "Unterminated comment!");
+      PrintError(TokStart, "unterminated comment");
       return true;
     case '*':
       // End of the comment?
@@ -543,7 +543,7 @@ tgtok::TokKind TGLexer::LexNumber() {
 
   // Requires at least one digit.
   if (CurPtr == NumStart)
-    return ReturnError(TokStart, "Invalid number");
+    return ReturnError(TokStart, "invalid number");
 
   errno = 0;
   if (IsMinus)
@@ -552,9 +552,9 @@ tgtok::TokKind TGLexer::LexNumber() {
     CurIntVal = strtoull(NumStart, nullptr, Base);
 
   if (errno == EINVAL)
-    return ReturnError(TokStart, "Invalid number");
+    return ReturnError(TokStart, "invalid number");
   if (errno == ERANGE)
-    return ReturnError(TokStart, "Number out of range");
+    return ReturnError(TokStart, "number out of range");
 
   return Base == 2 ? tgtok::BinaryIntVal : tgtok::IntVal;
 }
@@ -580,13 +580,13 @@ tgtok::TokKind TGLexer::LexBracket() {
     }
   }
 
-  return ReturnError(CodeStart - 2, "Unterminated code block");
+  return ReturnError(CodeStart - 2, "unterminated code block");
 }
 
 /// LexExclaim - Lex '!' and '![a-zA-Z]+'.
 tgtok::TokKind TGLexer::LexExclaim() {
   if (!isAlpha(*CurPtr))
-    return ReturnError(CurPtr - 1, "Invalid \"!operator\"");
+    return ReturnError(CurPtr - 1, "invalid \"!operator\"");
 
   const char *Start = CurPtr++;
   while (isAlpha(*CurPtr))
@@ -648,7 +648,8 @@ tgtok::TokKind TGLexer::LexExclaim() {
           .Case("repr", tgtok::XRepr)
           .Default(tgtok::Error);
 
-  return Kind != tgtok::Error ? Kind : ReturnError(Start-1, "Unknown operator");
+  return Kind != tgtok::Error ? Kind
+                              : ReturnError(Start - 1, "unknown operator");
 }
 
 bool TGLexer::prepExitInclude(bool IncludeStackMustBeEmpty) {
@@ -662,17 +663,17 @@ bool TGLexer::prepExitInclude(bool IncludeStackMustBeEmpty) {
 
   // Pop the preprocessing controls from the include stack.
   if (PrepIncludeStack.empty()) {
-    PrintFatalError("Preprocessor include stack is empty");
+    PrintFatalError("preprocessor include stack is empty");
   }
 
   PrepIncludeStack.pop_back();
 
   if (IncludeStackMustBeEmpty) {
     if (!PrepIncludeStack.empty())
-      PrintFatalError("Preprocessor include stack is not empty");
+      PrintFatalError("preprocessor include stack is not empty");
   } else {
     if (PrepIncludeStack.empty())
-      PrintFatalError("Preprocessor include stack is empty");
+      PrintFatalError("preprocessor include stack is empty");
   }
 
   return true;
@@ -732,7 +733,7 @@ bool TGLexer::prepEatPreprocessorDirective(tgtok::TokKind Kind) {
       return true;
     }
 
-  PrintFatalError("Unsupported preprocessing token in "
+  PrintFatalError("unsupported preprocessing token in "
                   "prepEatPreprocessorDirective()");
   return false;
 }
@@ -748,7 +749,7 @@ tgtok::TokKind TGLexer::lexPreprocessor(tgtok::TokKind Kind,
     StringRef MacroName = prepLexMacroName();
     StringRef IfTokName = Kind == tgtok::Ifdef ? "#ifdef" : "#ifndef";
     if (MacroName.empty())
-      return ReturnError(TokStart, "Expected macro name after " + IfTokName);
+      return ReturnError(TokStart, "expected macro name after " + IfTokName);
 
     bool MacroIsDefined = DefinedMacros.count(MacroName) != 0;
 
@@ -763,7 +764,7 @@ tgtok::TokKind TGLexer::lexPreprocessor(tgtok::TokKind Kind,
         {tgtok::Ifdef, MacroIsDefined, SMLoc::getFromPointer(TokStart)});
 
     if (!prepSkipDirectiveEnd())
-      return ReturnError(CurPtr, "Only comments are supported after " +
+      return ReturnError(CurPtr, "only comments are supported after " +
                                      IfTokName + " NAME");
 
     // If we were not processing tokens before this #ifdef,
@@ -794,7 +795,7 @@ tgtok::TokKind TGLexer::lexPreprocessor(tgtok::TokKind Kind,
 
     if (IfdefEntry.Kind != tgtok::Ifdef) {
       PrintError(TokStart, "double #else");
-      return ReturnError(IfdefEntry.SrcPos, "Previous #else is here");
+      return ReturnError(IfdefEntry.SrcPos, "previous #else is here");
     }
 
     // Replace the corresponding #ifdef's control with its negation
@@ -804,7 +805,7 @@ tgtok::TokKind TGLexer::lexPreprocessor(tgtok::TokKind Kind,
         {Kind, !IfdefEntry.IsDefined, SMLoc::getFromPointer(TokStart)});
 
     if (!prepSkipDirectiveEnd())
-      return ReturnError(CurPtr, "Only comments are supported after #else");
+      return ReturnError(CurPtr, "only comments are supported after #else");
 
     // If we were processing tokens before this #else,
     // we have to start skipping lines until the matching #endif.
@@ -827,12 +828,12 @@ tgtok::TokKind TGLexer::lexPreprocessor(tgtok::TokKind Kind,
 
     if (IfdefOrElseEntry.Kind != tgtok::Ifdef &&
         IfdefOrElseEntry.Kind != tgtok::Else) {
-      PrintFatalError("Invalid preprocessor control on the stack");
+      PrintFatalError("invalid preprocessor control on the stack");
       return tgtok::Error;
     }
 
     if (!prepSkipDirectiveEnd())
-      return ReturnError(CurPtr, "Only comments are supported after #endif");
+      return ReturnError(CurPtr, "only comments are supported after #endif");
 
     PrepIncludeStack.back()->pop_back();
 
@@ -847,15 +848,15 @@ tgtok::TokKind TGLexer::lexPreprocessor(tgtok::TokKind Kind,
   } else if (Kind == tgtok::Define) {
     StringRef MacroName = prepLexMacroName();
     if (MacroName.empty())
-      return ReturnError(TokStart, "Expected macro name after #define");
+      return ReturnError(TokStart, "expected macro name after #define");
 
     if (!DefinedMacros.insert(MacroName).second)
       PrintWarning(getLoc(),
-                   "Duplicate definition of macro: " + Twine(MacroName));
+                   "duplicate definition of macro: " + Twine(MacroName));
 
     if (!prepSkipDirectiveEnd())
       return ReturnError(CurPtr,
-                         "Only comments are supported after #define NAME");
+                         "only comments are supported after #define NAME");
 
     if (!ReturnNextLiveToken) {
       PrintFatalError("#define must be ignored during the lines skipping");
@@ -865,13 +866,13 @@ tgtok::TokKind TGLexer::lexPreprocessor(tgtok::TokKind Kind,
     return LexToken();
   }
 
-  PrintFatalError("Preprocessing directive is not supported");
+  PrintFatalError("preprocessing directive is not supported");
   return tgtok::Error;
 }
 
 bool TGLexer::prepSkipRegion(bool MustNeverBeFalse) {
   if (!MustNeverBeFalse)
-    PrintFatalError("Invalid recursion.");
+    PrintFatalError("invalid recursion.");
 
   do {
     // Skip all symbols to the line end.
@@ -917,7 +918,7 @@ bool TGLexer::prepSkipRegion(bool MustNeverBeFalse) {
     // due to #else or #endif.
     if (prepIsProcessingEnabled()) {
       if (Kind != tgtok::Else && Kind != tgtok::Endif) {
-        PrintFatalError("Tokens processing was enabled by an unexpected "
+        PrintFatalError("tokens processing was enabled by an unexpected "
                         "preprocessing directive");
         return false;
       }
@@ -1032,7 +1033,7 @@ bool TGLexer::prepSkipDirectiveEnd() {
           return false;
       } else {
         TokStart = CurPtr;
-        PrintError(CurPtr, "Unexpected character");
+        PrintError(CurPtr, "unexpected character");
         return false;
       }
 
@@ -1067,8 +1068,8 @@ void TGLexer::prepReportPreprocessorStackError() {
                     "empty control stack");
 
   auto &PrepControl = PrepIncludeStack.back()->back();
-  PrintError(CurBuf.end(), "Reached EOF without matching #endif");
-  PrintError(PrepControl.SrcPos, "The latest preprocessor control is here");
+  PrintError(CurBuf.end(), "reached EOF without matching #endif");
+  PrintError(PrepControl.SrcPos, "the latest preprocessor control is here");
 
   TokStart = CurPtr;
 }
diff --git a/llvm/lib/Target/PowerPC/PPCISelLowering.cpp b/llvm/lib/Target/PowerPC/PPCISelLowering.cpp
index ab31898e262e7e..d8f3095ed7fb68 100644
--- a/llvm/lib/Target/PowerPC/PPCISelLowering.cpp
+++ b/llvm/lib/Target/PowerPC/PPCISelLowering.cpp
@@ -200,6 +200,11 @@ PPCTargetLowering::PPCTargetLowering(const PPCTargetMachine &TM,
 
   setOperationAction(ISD::UADDO, isPPC64 ? MVT::i64 : MVT::i32, Custom);
 
+  // On P10, the default lowering generates better code using the
+  // setbc instruction.
+  if (!Subtarget.hasP10Vector() && isPPC64)
+    setOperationAction(ISD::SSUBO, MVT::i32, Custom);
+
   // Match BITREVERSE to customized fast code sequence in the td file.
   setOperationAction(ISD::BITREVERSE, MVT::i32, Legal);
   setOperationAction(ISD::BITREVERSE, MVT::i64, Legal);
@@ -12016,6 +12021,36 @@ SDValue PPCTargetLowering::LowerUaddo(SDValue Op, SelectionDAG &DAG) const {
   return Res;
 }
 
+SDValue PPCTargetLowering::LowerSSUBO(SDValue Op, SelectionDAG &DAG) const {
+
+  SDLoc dl(Op);
+
+  SDValue LHS64 = DAG.getNode(ISD::ANY_EXTEND, dl, MVT::i64, Op.getOperand(0));
+  SDValue RHS64 = DAG.getNode(ISD::ANY_EXTEND, dl, MVT::i64, Op.getOperand(1));
+
+  SDValue Sub = DAG.getNode(ISD::SUB, dl, MVT::i64, LHS64, RHS64);
+
+  SDValue Extsw = DAG.getNode(ISD::SIGN_EXTEND_INREG, dl, MVT::i64, Sub,
+                              DAG.getValueType(MVT::i32));
+
+  SDValue Xor = DAG.getNode(ISD::XOR, dl, MVT::i64, Extsw, Sub);
+
+  SDValue Addic = DAG.getNode(ISD::ADDC, dl, DAG.getVTList(MVT::i64, MVT::Glue),
+                              Xor, DAG.getConstant(-1, dl, MVT::i64));
+
+  SDValue Overflow =
+      DAG.getNode(ISD::SUBE, dl, DAG.getVTList(MVT::i64, MVT::Glue), Xor, Addic,
+                  Addic.getValue(1));
+
+  SDValue OverflowTrunc =
+      DAG.getNode(ISD::TRUNCATE, dl, Op.getNode()->getValueType(1), Overflow);
+  SDValue SubTrunc =
+      (Sub->getValueType(0) != Op.getNode()->getValueType(0))
+          ? DAG.getNode(ISD::TRUNCATE, dl, Op.getNode()->getValueType(0), Sub)
+          : Sub;
+  return DAG.getMergeValues({SubTrunc, OverflowTrunc}, dl);
+}
+
 /// LowerOperation - Provide custom lowering hooks for some operations.
 ///
 SDValue PPCTargetLowering::LowerOperation(SDValue Op, SelectionDAG &DAG) const {
@@ -12038,6 +12073,8 @@ SDValue PPCTargetLowering::LowerOperation(SDValue Op, SelectionDAG &DAG) const {
   case ISD::SETCC:              return LowerSETCC(Op, DAG);
   case ISD::INIT_TRAMPOLINE:    return LowerINIT_TRAMPOLINE(Op, DAG);
   case ISD::ADJUST_TRAMPOLINE:  return LowerADJUST_TRAMPOLINE(Op, DAG);
+  case ISD::SSUBO:
+    return LowerSSUBO(Op, DAG);
 
   case ISD::INLINEASM:
   case ISD::INLINEASM_BR:       return LowerINLINEASM(Op, DAG);
diff --git a/llvm/lib/Target/PowerPC/PPCISelLowering.h b/llvm/lib/Target/PowerPC/PPCISelLowering.h
index 0adbad86845973..dde45e4cf6f4ae 100644
--- a/llvm/lib/Target/PowerPC/PPCISelLowering.h
+++ b/llvm/lib/Target/PowerPC/PPCISelLowering.h
@@ -1279,6 +1279,7 @@ namespace llvm {
     SDValue LowerJumpTable(SDValue Op, SelectionDAG &DAG) const;
     SDValue LowerUaddo(SDValue Op, SelectionDAG &DAG) const;
     SDValue LowerSETCC(SDValue Op, SelectionDAG &DAG) const;
+    SDValue LowerSSUBO(SDValue Op, SelectionDAG &DAG) const;
     SDValue LowerINIT_TRAMPOLINE(SDValue Op, SelectionDAG &DAG) const;
     SDValue LowerADJUST_TRAMPOLINE(SDValue Op, SelectionDAG &DAG) const;
     SDValue LowerINLINEASM(SDValue Op, SelectionDAG &DAG) const;
diff --git a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
index d9377fe4b91a1a..11ed7d660be09e 100644
--- a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
+++ b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp
@@ -2547,6 +2547,17 @@ bool SPIRVInstructionSelector::selectIntrinsic(Register ResVReg,
     return selectExtInst(ResVReg, ResType, I, CL::rsqrt, GL::InverseSqrt);
   case Intrinsic::spv_sign:
     return selectSign(ResVReg, ResType, I);
+  case Intrinsic::spv_group_memory_barrier_with_group_sync: {
+    Register MemSemReg =
+        buildI32Constant(SPIRV::MemorySemantics::SequentiallyConsistent, I);
+    Register ScopeReg = buildI32Constant(SPIRV::Scope::Workgroup, I);
+    MachineBasicBlock &BB = *I.getParent();
+    return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpControlBarrier))
+        .addUse(ScopeReg)
+        .addUse(ScopeReg)
+        .addUse(MemSemReg)
+        .constrainAllUses(TII, TRI, RBI);
+  } break;
   case Intrinsic::spv_lifetime_start:
   case Intrinsic::spv_lifetime_end: {
     unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
diff --git a/llvm/test/Assembler/invalid-inttype.ll b/llvm/test/Assembler/invalid-inttype.ll
index c8aa7c66b79e4d..9e3c31148af2d6 100644
--- a/llvm/test/Assembler/invalid-inttype.ll
+++ b/llvm/test/Assembler/invalid-inttype.ll
@@ -1,5 +1,5 @@
 ; RUN: not llvm-as --disable-output %s 2>&1 | FileCheck -DFILE=%s %s
 
 ; i8388609 is the smallest integer type that can't be represented in LLVM IR
-; CHECK: [[FILE]]:[[@LINE+1]]:21: error: bitwidth for integer type out of range!
+; CHECK: [[FILE]]:[[@LINE+1]]:21: error: bitwidth for integer type out of range
 @i2 = common global i8388609 0, align 4
diff --git a/llvm/test/Assembler/invalid-name.ll b/llvm/test/Assembler/invalid-name.ll
index 74133e60df54d5..52e2bda3adbabd 100644
Binary files a/llvm/test/Assembler/invalid-name.ll and b/llvm/test/Assembler/invalid-name.ll differ
diff --git a/llvm/test/Assembler/invalid-name2.ll b/llvm/test/Assembler/invalid-name2.ll
index 8a848798a54caf..78da4dc3d1b8d0 100644
Binary files a/llvm/test/Assembler/invalid-name2.ll and b/llvm/test/Assembler/invalid-name2.ll differ
diff --git a/llvm/test/CodeGen/PowerPC/saddo-ssubo.ll b/llvm/test/CodeGen/PowerPC/saddo-ssubo.ll
index fd5f26ba35742f..7147257d27c4b8 100644
--- a/llvm/test/CodeGen/PowerPC/saddo-ssubo.ll
+++ b/llvm/test/CodeGen/PowerPC/saddo-ssubo.ll
@@ -129,12 +129,11 @@ entry:
 define i1 @test_ssubo_i32(i32 %a, i32 %b) nounwind {
 ; CHECK-LABEL: test_ssubo_i32:
 ; CHECK:       # %bb.0: # %entry
-; CHECK-NEXT:    sub 5, 3, 4
-; CHECK-NEXT:    cmpwi 1, 4, 0
-; CHECK-NEXT:    cmpw 5, 3
-; CHECK-NEXT:    li 3, 1
-; CHECK-NEXT:    creqv 20, 5, 0
-; CHECK-NEXT:    isel 3, 0, 3, 20
+; CHECK-NEXT:    sub 3, 3, 4
+; CHECK-NEXT:    extsw 4, 3
+; CHECK-NEXT:    xor 3, 4, 3
+; CHECK-NEXT:    addic 4, 3, -1
+; CHECK-NEXT:    subfe 3, 4, 3
 ; CHECK-NEXT:    blr
 entry:
   %res = call { i32, i1 } @llvm.ssub.with.overflow.i32(i32 %a, i32 %b) nounwind
diff --git a/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/group_memory_barrier_with_group_sync.ll b/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/group_memory_barrier_with_group_sync.ll
new file mode 100644
index 00000000000000..6955411a0e4e99
--- /dev/null
+++ b/llvm/test/CodeGen/SPIRV/hlsl-intrinsics/group_memory_barrier_with_group_sync.ll
@@ -0,0 +1,14 @@
+; RUN: llc -verify-machineinstrs -O0 -mtriple=spirv-unknown-unknown %s -o - | FileCheck %s
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+
+; CHECK: OpMemoryModel Logical GLSL450
+
+define void @test_group_memory_barrier_with_group_sync() {
+entry:
+  ; CHECK: %[[#TY:]] = OpTypeInt 32 0
+  ; CHECK-DAG: %[[#MEM_SEM:]] = OpConstant %[[#TY]] 16
+  ; CHECK-DAG: %[[#EXEC_AND_MEM_SCOPE:]] = OpConstant %[[#TY]] 2
+  ; CHECK: OpControlBarrier %[[#EXEC_AND_MEM_SCOPE]] %[[#EXEC_AND_MEM_SCOPE]] %[[#MEM_SEM]]
+  call void @llvm.spv.group.memory.barrier.with.group.sync()
+  ret void
+}
diff --git a/llvm/test/TableGen/64-bit-int.td b/llvm/test/TableGen/64-bit-int.td
index 2d2bdb8b560e2c..d2a2999c14e991 100644
--- a/llvm/test/TableGen/64-bit-int.td
+++ b/llvm/test/TableGen/64-bit-int.td
@@ -16,7 +16,7 @@ def {
 #ifdef OOR3
   bits<64> Val = 0x10000000000000000;
 #endif
-// CHECK-OOR: error: Number out of range
+// CHECK-OOR: error: number out of range
 
   bits<64> BinVal0 = 0x8000000000000000;
   bits<64> HexVal0 = 0b1000000000000000000000000000000000000000000000000000000000000000;
diff --git a/llvm/test/TableGen/invalid-macro-name-command-line.td b/llvm/test/TableGen/invalid-macro-name-command-line.td
index 0d2307997ebe54..7d19e8996639af 100644
--- a/llvm/test/TableGen/invalid-macro-name-command-line.td
+++ b/llvm/test/TableGen/invalid-macro-name-command-line.td
@@ -3,7 +3,7 @@
 // RUN: not llvm-tblgen %s -D_MAC# 2>&1 | FileCheck %s --check-prefix=CHECK-TEST-3
 // RUN: not llvm-tblgen %s -D 2>&1 | FileCheck %s --check-prefix=CHECK-TEST-4
 
-// CHECK-TEST-1: error: Invalid macro name `MACRO=1` specified on command line
-// CHECK-TEST-2: error: Invalid macro name `0MAC` specified on command line
-// CHECK-TEST-3: error: Invalid macro name `_MAC#` specified on command line
+// CHECK-TEST-1: error: invalid macro name `MACRO=1` specified on command line
+// CHECK-TEST-2: error: invalid macro name `0MAC` specified on command line
+// CHECK-TEST-3: error: invalid macro name `_MAC#` specified on command line
 // CHECK-TEST-4: for the -D option: requires a value!
diff --git a/llvm/test/TableGen/prep-diag1.td b/llvm/test/TableGen/prep-diag1.td
index 41b7d477c6942e..27f428f4fe9598 100644
--- a/llvm/test/TableGen/prep-diag1.td
+++ b/llvm/test/TableGen/prep-diag1.td
@@ -4,22 +4,22 @@
 // RUN: not llvm-tblgen -I %p %s 2>&1 | FileCheck --check-prefixes=DIAG3 %s
 
 #ifdef DIAG1
-// DIAG1: error: Only comments are supported after #define NAME
+// DIAG1: error: only comments are supported after #define NAME
 #define ENABLED1/*
 */class C;
 #endif // DIAG1
 
 #ifdef DIAG4
-// DIAG4: warning: Duplicate definition of macro: ENABLED1
+// DIAG4: warning: duplicate definition of macro: ENABLED1
 #define ENABLED1
 #define ENABLED1
 #endif // DIAG4
 
 #ifdef DIAG2
-// DIAG2: error: Only comments are supported after #ifdef NAME
+// DIAG2: error: only comments are supported after #ifdef NAME
 
 // Invalid #ifdef below should be detected even if DIAG2 is not defined.
-// DIAG3: error: Only comments are supported after #ifdef NAME
+// DIAG3: error: only comments are supported after #ifdef NAME
 #ifdef DIAG2/*
 */class C;
 #endif
diff --git a/llvm/test/TableGen/prep-diag10.td b/llvm/test/TableGen/prep-diag10.td
index eb387a07b066ca..cfcbab094ad73b 100644
--- a/llvm/test/TableGen/prep-diag10.td
+++ b/llvm/test/TableGen/prep-diag10.td
@@ -1,6 +1,6 @@
 // RUN: not llvm-tblgen -I %p %s 2>&1 | FileCheck %s
 
-// CHECK: error: Reached EOF without matching #endif
-// CHECK: error: The latest preprocessor control is here
+// CHECK: error: reached EOF without matching #endif
+// CHECK: error: the latest preprocessor control is here
 #ifdef DISABLED
 #else
diff --git a/llvm/test/TableGen/prep-diag11.td b/llvm/test/TableGen/prep-diag11.td
index 0042bc04f9e101..1fe8a8503076e5 100644
--- a/llvm/test/TableGen/prep-diag11.td
+++ b/llvm/test/TableGen/prep-diag11.td
@@ -1,7 +1,7 @@
 // RUN: not llvm-tblgen -I %p %s 2>&1 | FileCheck %s
 
-// CHECK: error: Reached EOF without matching #endif
-// CHECK: error: The latest preprocessor control is here
+// CHECK: error: reached EOF without matching #endif
+// CHECK: error: the latest preprocessor control is here
 #ifdef DISABLED
 #else
 #define ENABLED
diff --git a/llvm/test/TableGen/prep-diag12.td b/llvm/test/TableGen/prep-diag12.td
index c26301ee17ac2b..02ffa672b2fa05 100644
--- a/llvm/test/TableGen/prep-diag12.td
+++ b/llvm/test/TableGen/prep-diag12.td
@@ -1,7 +1,7 @@
 // RUN: not llvm-tblgen -I %p %s 2>&1 | FileCheck %s
 
-// CHECK: error: Reached EOF without matching #endif
-// CHECK: error: The latest preprocessor control is here
+// CHECK: error: reached EOF without matching #endif
+// CHECK: error: the latest preprocessor control is here
 #ifdef DISABLED
 #else
 #define ENABLED
diff --git a/llvm/test/TableGen/prep-diag13.td b/llvm/test/TableGen/prep-diag13.td
index aa3fdab4802d37..733a46a1618131 100644
--- a/llvm/test/TableGen/prep-diag13.td
+++ b/llvm/test/TableGen/prep-diag13.td
@@ -1,7 +1,7 @@
 // RUN: not llvm-tblgen -I %p %s 2>&1 | FileCheck %s
 
-// CHECK: error: Reached EOF without matching #endif
-// CHECK: error: The latest preprocessor control is here
+// CHECK: error: reached EOF without matching #endif
+// CHECK: error: the latest preprocessor control is here
 #ifdef DISABLED
 /*
 #else
diff --git a/llvm/test/TableGen/prep-diag14.td b/llvm/test/TableGen/prep-diag14.td
index cae9bc3b7f5b6c..a3216ee4f47125 100644
--- a/llvm/test/TableGen/prep-diag14.td
+++ b/llvm/test/TableGen/prep-diag14.td
@@ -1,6 +1,6 @@
 // RUN: not llvm-tblgen -I %p %s 2>&1 | FileCheck %s
 
-// CHECK: error: Reached EOF without matching #endif
-// CHECK: error: The latest preprocessor control is here
+// CHECK: error: reached EOF without matching #endif
+// CHECK: error: the latest preprocessor control is here
 #ifdef DISABLED
 // #endif
diff --git a/llvm/test/TableGen/prep-diag2.td b/llvm/test/TableGen/prep-diag2.td
index 741026b9c8a2d6..e51490600ff64f 100644
--- a/llvm/test/TableGen/prep-diag2.td
+++ b/llvm/test/TableGen/prep-diag2.td
@@ -2,10 +2,10 @@
 // RUN: not llvm-tblgen -I %p %s 2>&1 | FileCheck --check-prefixes=DIAG2 %s
 
 #ifdef DIAG1
-// DIAG1: error: Only comments are supported after #else
+// DIAG1: error: only comments are supported after #else
 
 // Invalid #else below should be detected even if DIAG1 is not defined.
-// DIAG2: error: Only comments are supported after #else
+// DIAG2: error: only comments are supported after #else
 #ifdef DIAG2//DIAG2
 #else/*
 */class C;
diff --git a/llvm/test/TableGen/prep-diag3.td b/llvm/test/TableGen/prep-diag3.td
index fbedfa290b9947..0b4d40307b40b0 100644
--- a/llvm/test/TableGen/prep-diag3.td
+++ b/llvm/test/TableGen/prep-diag3.td
@@ -2,10 +2,10 @@
 // RUN: not llvm-tblgen -I %p %s 2>&1 | FileCheck --check-prefixes=DIAG2 %s
 
 #ifdef DIAG1
-// DIAG1: error: Only comments are supported after #endif
+// DIAG1: error: only comments are supported after #endif
 
 // Invalid #else below should be detected even if DIAG1 is not defined.
-// DIAG2: error: Only comments are supported after #endif
+// DIAG2: error: only comments are supported after #endif
 #ifdef DIAG2//DIAG2
 #else/*!DIAG2*/
 #endif/* !DIAG2
diff --git a/llvm/test/TableGen/prep-diag4.td b/llvm/test/TableGen/prep-diag4.td
index 4661ef8667d23f..ead116ebde0de8 100644
--- a/llvm/test/TableGen/prep-diag4.td
+++ b/llvm/test/TableGen/prep-diag4.td
@@ -1,7 +1,7 @@
 // RUN: not llvm-tblgen -I %p %s 2>&1 | FileCheck %s
 
 // CHECK: error: double #else
-// CHECK: error: Previous #else is here
+// CHECK: error: previous #else is here
 #ifdef DIAG1
 #else
 #else
diff --git a/llvm/test/TableGen/prep-diag6.td b/llvm/test/TableGen/prep-diag6.td
index f4202d115da59a..bf1cd3d3490b5e 100644
--- a/llvm/test/TableGen/prep-diag6.td
+++ b/llvm/test/TableGen/prep-diag6.td
@@ -1,6 +1,6 @@
 // RUN: not llvm-tblgen -I %p %s 2>&1 | FileCheck %s
 
-// CHECK: error: Expected macro name after #ifdef
+// CHECK: error: expected macro name after #ifdef
 #ifdef
 #else
 #else
diff --git a/llvm/test/TableGen/prep-diag8.td b/llvm/test/TableGen/prep-diag8.td
index 7a7bde62c79c4e..82797d6cf4a62d 100644
--- a/llvm/test/TableGen/prep-diag8.td
+++ b/llvm/test/TableGen/prep-diag8.td
@@ -1,5 +1,5 @@
 // RUN: not llvm-tblgen -I %p %s 2>&1 | FileCheck %s
 
-// CHECK: error: Expected macro name after #define
+// CHECK: error: expected macro name after #define
 #define
 #endif
diff --git a/llvm/test/TableGen/prep-diag9.td b/llvm/test/TableGen/prep-diag9.td
index 4ecff575cdc7bb..6ad208104301bc 100644
--- a/llvm/test/TableGen/prep-diag9.td
+++ b/llvm/test/TableGen/prep-diag9.td
@@ -1,5 +1,5 @@
 // RUN: not llvm-tblgen -I %p %s 2>&1 | FileCheck %s
 
-// CHECK: error: Reached EOF without matching #endif
-// CHECK: error: The latest preprocessor control is here
+// CHECK: error: reached EOF without matching #endif
+// CHECK: error: the latest preprocessor control is here
 #ifdef DISABLED
diff --git a/llvm/test/TableGen/prep-ifndef-diag-1.td b/llvm/test/TableGen/prep-ifndef-diag-1.td
index 941f2d377a98a7..4a0d0754ed7906 100644
--- a/llvm/test/TableGen/prep-ifndef-diag-1.td
+++ b/llvm/test/TableGen/prep-ifndef-diag-1.td
@@ -1,4 +1,4 @@
 // RUN: not llvm-tblgen %s 2>&1 | FileCheck %s
 
-// CHECK: error: Expected macro name after #ifndef
+// CHECK: error: expected macro name after #ifndef
 #ifndef 1
diff --git a/llvm/test/TableGen/prep-ifndef-diag-2.td b/llvm/test/TableGen/prep-ifndef-diag-2.td
index 7b5f9dfd24b786..c89cbab08e5c5c 100644
--- a/llvm/test/TableGen/prep-ifndef-diag-2.td
+++ b/llvm/test/TableGen/prep-ifndef-diag-2.td
@@ -1,4 +1,4 @@
 // RUN: not llvm-tblgen %s 2>&1 | FileCheck %s
 
-// CHECK: error: Only comments are supported after #ifndef NAME
+// CHECK: error: only comments are supported after #ifndef NAME
 #ifndef MACRO 42
diff --git a/llvm/test/TableGen/unterminated-c-comment.td b/llvm/test/TableGen/unterminated-c-comment.td
index 0f4cd9d633c66d..b5b995342be744 100644
--- a/llvm/test/TableGen/unterminated-c-comment.td
+++ b/llvm/test/TableGen/unterminated-c-comment.td
@@ -1,5 +1,5 @@
 // RUN: not llvm-tblgen -I %p %s 2>&1 | FileCheck %s
 
-// CHECK: error: Unterminated comment!
+// CHECK: error: unterminated comment
 
 include "unterminated-c-comment-include.inc" */
diff --git a/llvm/test/TableGen/unterminated-code-block.td b/llvm/test/TableGen/unterminated-code-block.td
index d6b6f50827a672..5bd4cd7e17d827 100644
--- a/llvm/test/TableGen/unterminated-code-block.td
+++ b/llvm/test/TableGen/unterminated-code-block.td
@@ -1,5 +1,5 @@
 // RUN: not llvm-tblgen -I %p %s 2>&1 | FileCheck %s
 
-// CHECK: error: Unterminated code block
+// CHECK: error: unterminated code block
 
 include "unterminated-code-block-include.inc" }]>;