merge main into amd-staging

Change-Id: I2ce52449addec5130cfee26dc9de4b393e9773c3

.github/workflows/containers/github-action-ci/stage1.Dockerfile

@@ -2,7 +2,7 @@ FROM docker.io/library/ubuntu:22.04 as base
 ENV LLVM_SYSROOT=/opt/llvm
 
 FROM base as stage1-toolchain
-ENV LLVM_VERSION=18.1.8
+ENV LLVM_VERSION=19.1.2
 
 RUN apt-get update && \
     apt-get install -y \

clang/test/CodeGenCUDA/bf16.cu

@@ -25,7 +25,7 @@ __device__ void test_arg(__bf16 *out, __bf16 in) {
 __device__ __bf16 test_ret( __bf16 in) {
 // CHECK:        ld.param.b16    %[[R:rs[0-9]+]], [_Z8test_retDF16b_param_0];
   return in;
-// CHECK:        st.param.b16    [func_retval0+0], %[[R]]
+// CHECK:        st.param.b16    [func_retval0], %[[R]]
 // CHECK:        ret;
 }
 
@@ -35,15 +35,15 @@ __device__ __bf16 external_func( __bf16 in);
 // CHECK:        .param .align 2 .b8 _Z9test_callDF16b_param_0[2]
 __device__ __bf16 test_call( __bf16 in) {
 // CHECK:        ld.param.b16    %[[R:rs[0-9]+]], [_Z9test_callDF16b_param_0];
-// CHECK:        st.param.b16    [param0+0], %[[R]];
+// CHECK:        st.param.b16    [param0], %[[R]];
 // CHECK:        .param .align 2 .b8 retval0[2];
 // CHECK:        call.uni (retval0),
 // CHECK-NEXT:   _Z13external_funcDF16b,
 // CHECK-NEXT:   (
 // CHECK-NEXT:   param0
 // CHECK-NEXT    );
-// CHECK:        ld.param.b16    %[[RET:rs[0-9]+]], [retval0+0];
+// CHECK:        ld.param.b16    %[[RET:rs[0-9]+]], [retval0];
   return external_func(in);
-// CHECK:        st.param.b16    [func_retval0+0], %[[RET]]
+// CHECK:        st.param.b16    [func_retval0], %[[RET]]
 // CHECK:        ret;
 }

compiler-rt/lib/hwasan/CMakeLists.txt

@@ -24,16 +24,19 @@ foreach(arch ${HWASAN_SUPPORTED_ARCH})
   if(${arch} MATCHES "aarch64")
     list(APPEND HWASAN_RTL_SOURCES
       hwasan_setjmp_aarch64.S
-      hwasan_tag_mismatch_aarch64.S)
+      hwasan_tag_mismatch_aarch64.S
+      )
   endif()
   if(${arch} MATCHES "riscv64")
     list(APPEND HWASAN_RTL_SOURCES
       hwasan_setjmp_riscv64.S
-      hwasan_tag_mismatch_riscv64.S)
+      hwasan_tag_mismatch_riscv64.S
+      )
   endif()
   if(${arch} MATCHES "x86_64")
     list(APPEND HWASAN_RTL_SOURCES
-      hwasan_setjmp_x86_64.S)
+      hwasan_setjmp_x86_64.S
+      )
   endif()
 endforeach()
 

lldb/include/lldb/Core/Module.h

@@ -30,6 +30,7 @@
 
 #include "llvm/ADT/DenseSet.h"
 #include "llvm/ADT/STLFunctionalExtras.h"
+#include "llvm/ADT/StableHashing.h"
 #include "llvm/ADT/StringRef.h"
 #include "llvm/Support/Chrono.h"
 
@@ -1057,8 +1058,11 @@ class Module : public std::enable_shared_from_this<Module>,
   /// time for the symbol tables can be aggregated here.
   StatsDuration m_symtab_index_time;
 
-  std::once_flag m_optimization_warning;
-  std::once_flag m_language_warning;
+  /// A set of hashes of all warnings and errors, to avoid reporting them
+  /// multiple times to the same Debugger.
+  llvm::DenseMap<llvm::stable_hash, std::unique_ptr<std::once_flag>>
+      m_shown_diagnostics;
+  std::recursive_mutex m_diagnostic_mutex;
 
   void SymbolIndicesToSymbolContextList(Symtab *symtab,
                                         std::vector<uint32_t> &symbol_indexes,
@@ -1086,6 +1090,7 @@ class Module : public std::enable_shared_from_this<Module>,
   void ReportWarning(const llvm::formatv_object_base &payload);
   void ReportError(const llvm::formatv_object_base &payload);
   void ReportErrorIfModifyDetected(const llvm::formatv_object_base &payload);
+  std::once_flag *GetDiagnosticOnceFlag(llvm::StringRef msg);
 };
 
 } // namespace lldb_private

lldb/source/Core/Module.cpp

@@ -1093,8 +1093,8 @@ void Module::ReportWarningOptimization(
   ss << file_name
      << " was compiled with optimization - stepping may behave "
         "oddly; variables may not be available.";
-  Debugger::ReportWarning(std::string(ss.GetString()), debugger_id,
-                          &m_optimization_warning);
+  llvm::StringRef msg = ss.GetString();
+  Debugger::ReportWarning(msg.str(), debugger_id, GetDiagnosticOnceFlag(msg));
 }
 
 void Module::ReportWarningUnsupportedLanguage(
@@ -1104,8 +1104,8 @@ void Module::ReportWarningUnsupportedLanguage(
      << Language::GetNameForLanguageType(language)
      << "\". "
         "Inspection of frame variables will be limited.";
-  Debugger::ReportWarning(std::string(ss.GetString()), debugger_id,
-                          &m_language_warning);
+  llvm::StringRef msg = ss.GetString();
+  Debugger::ReportWarning(msg.str(), debugger_id, GetDiagnosticOnceFlag(msg));
 }
 
 void Module::ReportErrorIfModifyDetected(
@@ -1125,20 +1125,29 @@ void Module::ReportErrorIfModifyDetected(
   }
 }
 
+std::once_flag *Module::GetDiagnosticOnceFlag(llvm::StringRef msg) {
+  std::lock_guard<std::recursive_mutex> guard(m_diagnostic_mutex);
+  auto &once_ptr = m_shown_diagnostics[llvm::stable_hash_name(msg)];
+  if (!once_ptr)
+    once_ptr = std::make_unique<std::once_flag>();
+  return once_ptr.get();
+}
+
 void Module::ReportError(const llvm::formatv_object_base &payload) {
   StreamString strm;
   GetDescription(strm.AsRawOstream(), lldb::eDescriptionLevelBrief);
-  strm.PutChar(' ');
-  strm.PutCString(payload.str());
-  Debugger::ReportError(strm.GetString().str());
+  std::string msg = payload.str();
+  strm << ' ' << msg;
+  Debugger::ReportError(strm.GetString().str(), {}, GetDiagnosticOnceFlag(msg));
 }
 
 void Module::ReportWarning(const llvm::formatv_object_base &payload) {
   StreamString strm;
   GetDescription(strm.AsRawOstream(), lldb::eDescriptionLevelFull);
-  strm.PutChar(' ');
-  strm.PutCString(payload.str());
-  Debugger::ReportWarning(std::string(strm.GetString()));
+  std::string msg = payload.str();
+  strm << ' ' << msg;
+  Debugger::ReportWarning(strm.GetString().str(), {},
+                          GetDiagnosticOnceFlag(msg));
 }
 
 void Module::LogMessage(Log *log, const llvm::formatv_object_base &payload) {

lldb/source/Plugins/SymbolFile/DWARF/SymbolFileDWARF.cpp

@@ -2069,13 +2069,15 @@ void SymbolFileDWARF::UpdateExternalModuleListIfNeeded() {
     Status error = ModuleList::GetSharedModule(dwo_module_spec, module_sp,
                                                nullptr, nullptr, nullptr);
     if (!module_sp) {
+      // ReportWarning also rate-limits based on the warning string,
+      // but in a -gmodules build, each object file has a similar DAG
+      // of module dependencies that would all be listed here.
       GetObjectFile()->GetModule()->ReportWarning(
-          "{0:x16}: unable to locate module needed for external types: "
-          "{1}\nerror: {2}\nDebugging will be degraded due to missing "
-          "types. Rebuilding the project will regenerate the needed "
-          "module files.",
-          die.GetOffset(), dwo_module_spec.GetFileSpec().GetPath().c_str(),
-          error.AsCString("unknown error"));
+          "{0}", error.AsCString("unknown error"));
+      GetObjectFile()->GetModule()->ReportWarning(
+          "Unable to locate module needed for external types.\n"
+          "Debugging will be degraded due to missing types. Rebuilding the "
+          "project will regenerate the needed module files.");
       continue;
     }
 
@@ -2095,12 +2097,11 @@ void SymbolFileDWARF::UpdateExternalModuleListIfNeeded() {
 
     if (dwo_id != dwo_dwo_id) {
       GetObjectFile()->GetModule()->ReportWarning(
-          "{0:x16}: Module {1} is out-of-date (hash mismatch). Type "
-          "information "
-          "from this module may be incomplete or inconsistent with the rest of "
-          "the program. Rebuilding the project will regenerate the needed "
-          "module files.",
-          die.GetOffset(), dwo_module_spec.GetFileSpec().GetPath().c_str());
+          "Module {0} is out-of-date (hash mismatch).\n"
+          "Type information from this module may be incomplete or inconsistent "
+          "with the rest of the program. Rebuilding the project will "
+          "regenerate the needed module files.",
+          dwo_module_spec.GetFileSpec().GetPath());
     }
   }
 }

lldb/test/Shell/SymbolFile/DWARF/TestDedupWarnings.test

@@ -0,0 +1,22 @@
+# REQUIRES: system-darwin
+# Test the rate-limiting of module not found warnings.
+# RUN: rm -rf %t
+# RUN: mkdir -p %t
+
+# RUN: echo 'module "C" { header "c.h" }' >%t/module.modulemap
+# RUN: echo 'struct c {};' >>%t/c.h
+# RUN: echo '@import C;'                  >%t/a.m
+# RUN: echo 'struct a { struct c c; } a;' >>%t/a.m
+# RUN: echo '@import C;'                  >%t/b.m
+# RUN: echo 'struct b { struct c c; } b;' >>%t/b.m
+# RUN: echo 'int main() {}'               >>%t/b.m
+
+# RUN: %clang_host -fmodules -Xclang -fmodules-cache-path=%t/cache -I%t -g -gmodules %t/a.m -o %t/a.o -c
+# RUN: %clang_host -fmodules -Xclang -fmodules-cache-path=%t/cache -I%t -g -gmodules %t/b.m -o %t/b.o -c
+# RUN: %clang_host %t/a.o %t/b.o -o %t/a.out
+# RUN: rm -rf %t/cache
+# RUN: %lldb %t/a.out -o "b main" -o run -o "p a" -o "p b" -o q 2>&1 | FileCheck %s
+# CHECK: {{[ab]}}.o{{.*}}/cache/{{.*}}/C-{{.*}}.pcm' does not exist
+# CHECK-NOT: /cache/{{.*}}/C-{.*}.pcm' does not exist
+# CHECK: {{[ab]}}.o{{.*}}/cache/{{.*}}/C-{{.*}}.pcm' does not exist
+# CHECK-NOT: /cache/{{.*}}/C-{.*}.pcm' does not exist

llvm/include/llvm/CodeGen/TargetLowering.h

@@ -3223,6 +3223,9 @@ class TargetLoweringBase {
   /// not legal, but should return true if those types will eventually legalize
   /// to types that support FMAs. After legalization, it will only be called on
   /// types that support FMAs (via Legal or Custom actions)
+  ///
+  /// Targets that care about soft float support should return false when soft
+  /// float code is being generated (i.e. use-soft-float).
   virtual bool isFMAFasterThanFMulAndFAdd(const MachineFunction &MF,
                                           EVT) const {
     return false;

llvm/lib/Target/ARM/ARMISelLowering.cpp

@@ -19354,6 +19354,9 @@ bool ARMTargetLowering::allowTruncateForTailCall(Type *Ty1, Type *Ty2) const {
 /// patterns (and we don't have the non-fused floating point instruction).
 bool ARMTargetLowering::isFMAFasterThanFMulAndFAdd(const MachineFunction &MF,
                                                    EVT VT) const {
+  if (Subtarget->useSoftFloat())
+    return false;
+
   if (!VT.isSimple())
     return false;
 

llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.cpp

@@ -363,6 +363,16 @@ void NVPTXInstPrinter::printMemOperand(const MCInst *MI, int OpNum,
   }
 }
 
+void NVPTXInstPrinter::printOffseti32imm(const MCInst *MI, int OpNum,
+                                         raw_ostream &O, const char *Modifier) {
+  auto &Op = MI->getOperand(OpNum);
+  assert(Op.isImm() && "Invalid operand");
+  if (Op.getImm() != 0) {
+    O << "+";
+    printOperand(MI, OpNum, O);
+  }
+}
+
 void NVPTXInstPrinter::printProtoIdent(const MCInst *MI, int OpNum,
                                        raw_ostream &O, const char *Modifier) {
   const MCOperand &Op = MI->getOperand(OpNum);

llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXInstPrinter.h

@@ -45,6 +45,8 @@ class NVPTXInstPrinter : public MCInstPrinter {
                     const char *Modifier = nullptr);
   void printMemOperand(const MCInst *MI, int OpNum,
                        raw_ostream &O, const char *Modifier = nullptr);
+  void printOffseti32imm(const MCInst *MI, int OpNum, raw_ostream &O,
+                         const char *Modifier = nullptr);
   void printProtoIdent(const MCInst *MI, int OpNum,
                        raw_ostream &O, const char *Modifier = nullptr);
   void printPrmtMode(const MCInst *MI, int OpNum, raw_ostream &O,