From 7154e7518563e5d12050f3082d32753ec9cea438 Mon Sep 17 00:00:00 2001 From: Razvan Lupusoru Date: Wed, 11 Feb 2026 11:13:04 -0800 Subject: [PATCH 001/705] [mlir][acc] Add gpu module and data layout utilities (#180988) Add utility functions to support OpenACC code generation: - getOrCreateGPUModule for GPU module management with configurable naming - getDataLayout for querying data layout from operations with explicit spec or module fallback --- .../Dialect/OpenACC/Analysis/OpenACCSupport.h | 37 ++++++++ .../mlir/Dialect/OpenACC/OpenACCUtilsCG.h | 40 +++++++++ .../mlir/Dialect/OpenACC/OpenACCUtilsGPU.h | 45 ++++++++++ .../Dialect/OpenACC/Analysis/CMakeLists.txt | 1 + .../OpenACC/Analysis/OpenACCSupport.cpp | 9 ++ mlir/lib/Dialect/OpenACC/Utils/CMakeLists.txt | 4 + .../Dialect/OpenACC/Utils/OpenACCUtilsCG.cpp | 55 ++++++++++++ .../Dialect/OpenACC/Utils/OpenACCUtilsGPU.cpp | 47 ++++++++++ mlir/unittests/Dialect/OpenACC/CMakeLists.txt | 3 + .../Dialect/OpenACC/OpenACCUtilsCGTest.cpp | 76 ++++++++++++++++ .../Dialect/OpenACC/OpenACCUtilsGPUTest.cpp | 89 +++++++++++++++++++ 11 files changed, 406 insertions(+) create mode 100644 mlir/include/mlir/Dialect/OpenACC/OpenACCUtilsCG.h create mode 100644 mlir/include/mlir/Dialect/OpenACC/OpenACCUtilsGPU.h create mode 100644 mlir/lib/Dialect/OpenACC/Utils/OpenACCUtilsCG.cpp create mode 100644 mlir/lib/Dialect/OpenACC/Utils/OpenACCUtilsGPU.cpp create mode 100644 mlir/unittests/Dialect/OpenACC/OpenACCUtilsCGTest.cpp create mode 100644 mlir/unittests/Dialect/OpenACC/OpenACCUtilsGPUTest.cpp diff --git a/mlir/include/mlir/Dialect/OpenACC/Analysis/OpenACCSupport.h b/mlir/include/mlir/Dialect/OpenACC/Analysis/OpenACCSupport.h index 31ef05c2f1743..8649160aa46d7 100644 --- a/mlir/include/mlir/Dialect/OpenACC/Analysis/OpenACCSupport.h +++ b/mlir/include/mlir/Dialect/OpenACC/Analysis/OpenACCSupport.h @@ -50,7 +50,9 @@ #ifndef MLIR_DIALECT_OPENACC_ANALYSIS_OPENACCSUPPORT_H #define MLIR_DIALECT_OPENACC_ANALYSIS_OPENACCSUPPORT_H +#include "mlir/Dialect/GPU/IR/GPUDialect.h" #include "mlir/Dialect/OpenACC/OpenACCUtils.h" +#include "mlir/Dialect/OpenACC/OpenACCUtilsGPU.h" #include "mlir/IR/Remarks.h" #include "mlir/IR/Value.h" #include "mlir/Pass/AnalysisManager.h" @@ -95,6 +97,10 @@ struct OpenACCSupportTraits { /// Check if a value use is legal in an OpenACC region. virtual bool isValidValueUse(Value v, mlir::Region ®ion) = 0; + + /// Get or optionally create a GPU module in the given module. + virtual std::optional + getOrCreateGPUModule(ModuleOp mod, bool create, llvm::StringRef name) = 0; }; /// SFINAE helpers to detect if implementation has optional methods @@ -125,6 +131,16 @@ struct OpenACCSupportTraits { llvm::is_detected, llvm::StringRef>; + template + using getOrCreateGPUModule_t = + decltype(std::declval().getOrCreateGPUModule( + std::declval()...)); + + template + using has_getOrCreateGPUModule = + llvm::is_detected; + /// This class wraps a concrete OpenACCSupport implementation and forwards /// interface calls to it. This provides type erasure, allowing different /// implementation types to be used interchangeably without inheritance. @@ -172,6 +188,15 @@ struct OpenACCSupportTraits { return acc::isValidValueUse(v, region); } + std::optional + getOrCreateGPUModule(ModuleOp mod, bool create, + llvm::StringRef name) final { + if constexpr (has_getOrCreateGPUModule::value) + return impl.getOrCreateGPUModule(mod, create, name); + else + return acc::getOrCreateGPUModule(mod, create, name); + } + private: ImplT impl; }; @@ -270,6 +295,18 @@ class OpenACCSupport { /// \param region The MLIR region in which the legality is checked. bool isValidValueUse(Value v, Region ®ion); + /// Get or optionally create a GPU module in the given module. + /// + /// \param mod The module to search or create the GPU module in. + /// \param create If true (default), create the GPU module if it doesn't + /// exist. + /// \param name The name for the GPU module. If empty, implementation uses its + /// default name. + /// \return The GPU module if found or created, std::nullopt otherwise. + std::optional + getOrCreateGPUModule(ModuleOp mod, bool create = true, + llvm::StringRef name = ""); + /// Signal that this analysis should always be preserved so that /// underlying implementation registration is not lost. bool isInvalidated(const AnalysisManager::PreservedAnalyses &pa) { diff --git a/mlir/include/mlir/Dialect/OpenACC/OpenACCUtilsCG.h b/mlir/include/mlir/Dialect/OpenACC/OpenACCUtilsCG.h new file mode 100644 index 0000000000000..7bead720b1077 --- /dev/null +++ b/mlir/include/mlir/Dialect/OpenACC/OpenACCUtilsCG.h @@ -0,0 +1,40 @@ +//===- OpenACCUtilsCG.h - OpenACC Code Generation Utilities -----*- C++ -*-===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// +// This file defines utility functions for OpenACC code generation, including +// data layout and type-related utilities. +// +//===----------------------------------------------------------------------===// + +#ifndef MLIR_DIALECT_OPENACC_OPENACCUTILSCG_H_ +#define MLIR_DIALECT_OPENACC_OPENACCUTILSCG_H_ + +#include "mlir/Interfaces/DataLayoutInterfaces.h" +#include + +namespace mlir { +class Operation; + +namespace acc { + +/// Get the data layout for an operation. +/// +/// Attempts to get the data layout from the operation or its parent module. +/// If `allowDefault` is true (default), a default data layout may be +/// constructed when no explicit data layout spec is found. +/// +/// \param op The operation to get the data layout for. +/// \param allowDefault If true, allow returning a default data layout. +/// \return The data layout if available, std::nullopt otherwise. +std::optional getDataLayout(Operation *op, + bool allowDefault = true); + +} // namespace acc +} // namespace mlir + +#endif // MLIR_DIALECT_OPENACC_OPENACCUTILSCG_H_ diff --git a/mlir/include/mlir/Dialect/OpenACC/OpenACCUtilsGPU.h b/mlir/include/mlir/Dialect/OpenACC/OpenACCUtilsGPU.h new file mode 100644 index 0000000000000..aebd15daca060 --- /dev/null +++ b/mlir/include/mlir/Dialect/OpenACC/OpenACCUtilsGPU.h @@ -0,0 +1,45 @@ +//===- OpenACCUtilsGPU.h - OpenACC GPU Utilities -----------------*- C++-*-===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// +// This file defines utility functions for OpenACC that depend on the GPU +// dialect. +// +//===----------------------------------------------------------------------===// + +#ifndef MLIR_DIALECT_OPENACC_OPENACCUTILSGPU_H_ +#define MLIR_DIALECT_OPENACC_OPENACCUTILSGPU_H_ + +#include "mlir/Dialect/GPU/IR/GPUDialect.h" +#include "mlir/IR/BuiltinOps.h" +#include + +namespace mlir { +namespace acc { + +/// Default GPU module name used by OpenACC. +constexpr llvm::StringLiteral kDefaultGPUModuleName = "acc_gpu_module"; + +/// Get or create a GPU module in the given module. +/// +/// If a GPU module with the specified name already exists, it is returned. +/// If `create` is true and no GPU module exists, one is created. +/// If `create` is false and no GPU module exists, std::nullopt is returned. +/// +/// \param mod The module to search or create the GPU module in. +/// \param create If true (default), create the GPU module if it doesn't exist. +/// \param name The name for the GPU module. If empty, uses +/// kDefaultGPUModuleName. +/// \return The GPU module if found or created, std::nullopt otherwise. +std::optional +getOrCreateGPUModule(ModuleOp mod, bool create = true, + llvm::StringRef name = kDefaultGPUModuleName); + +} // namespace acc +} // namespace mlir + +#endif // MLIR_DIALECT_OPENACC_OPENACCUTILSGPU_H_ diff --git a/mlir/lib/Dialect/OpenACC/Analysis/CMakeLists.txt b/mlir/lib/Dialect/OpenACC/Analysis/CMakeLists.txt index f305068e1b3bc..f2c13608855b9 100644 --- a/mlir/lib/Dialect/OpenACC/Analysis/CMakeLists.txt +++ b/mlir/lib/Dialect/OpenACC/Analysis/CMakeLists.txt @@ -5,6 +5,7 @@ add_mlir_dialect_library(MLIROpenACCAnalysis ${MLIR_MAIN_INCLUDE_DIR}/mlir/Dialect/OpenACC LINK_LIBS PUBLIC + MLIRGPUDialect MLIRIR MLIROpenACCDialect MLIROpenACCUtils diff --git a/mlir/lib/Dialect/OpenACC/Analysis/OpenACCSupport.cpp b/mlir/lib/Dialect/OpenACC/Analysis/OpenACCSupport.cpp index f204ace4609ab..3875122152194 100644 --- a/mlir/lib/Dialect/OpenACC/Analysis/OpenACCSupport.cpp +++ b/mlir/lib/Dialect/OpenACC/Analysis/OpenACCSupport.cpp @@ -12,6 +12,7 @@ #include "mlir/Dialect/OpenACC/Analysis/OpenACCSupport.h" #include "mlir/Dialect/OpenACC/OpenACCUtils.h" +#include "mlir/Dialect/OpenACC/OpenACCUtilsGPU.h" namespace mlir { namespace acc { @@ -63,5 +64,13 @@ bool OpenACCSupport::isValidValueUse(Value v, Region ®ion) { return acc::isValidValueUse(v, region); } +std::optional +OpenACCSupport::getOrCreateGPUModule(ModuleOp mod, bool create, + llvm::StringRef name) { + if (impl) + return impl->getOrCreateGPUModule(mod, create, name); + return acc::getOrCreateGPUModule(mod, create, name); +} + } // namespace acc } // namespace mlir diff --git a/mlir/lib/Dialect/OpenACC/Utils/CMakeLists.txt b/mlir/lib/Dialect/OpenACC/Utils/CMakeLists.txt index 532ba90355b44..baf248fd9b2ae 100644 --- a/mlir/lib/Dialect/OpenACC/Utils/CMakeLists.txt +++ b/mlir/lib/Dialect/OpenACC/Utils/CMakeLists.txt @@ -1,5 +1,7 @@ add_mlir_dialect_library(MLIROpenACCUtils OpenACCUtils.cpp + OpenACCUtilsCG.cpp + OpenACCUtilsGPU.cpp OpenACCUtilsLoop.cpp OpenACCUtilsTiling.cpp @@ -18,6 +20,8 @@ add_mlir_dialect_library(MLIROpenACCUtils LINK_LIBS PUBLIC MLIRArithDialect MLIRArithUtils + MLIRDataLayoutInterfaces + MLIRGPUDialect MLIROpenACCDialect MLIRIR MLIRSCFDialect diff --git a/mlir/lib/Dialect/OpenACC/Utils/OpenACCUtilsCG.cpp b/mlir/lib/Dialect/OpenACC/Utils/OpenACCUtilsCG.cpp new file mode 100644 index 0000000000000..5c5c453f2cae0 --- /dev/null +++ b/mlir/lib/Dialect/OpenACC/Utils/OpenACCUtilsCG.cpp @@ -0,0 +1,55 @@ +//===- OpenACCUtilsCG.cpp - OpenACC Code Generation Utilities -------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// +// This file implements utility functions for OpenACC code generation. +// +//===----------------------------------------------------------------------===// + +#include "mlir/Dialect/OpenACC/OpenACCUtilsCG.h" +#include "mlir/IR/BuiltinOps.h" + +namespace mlir { +namespace acc { + +std::optional getDataLayout(Operation *op, bool allowDefault) { + if (!op) + return std::nullopt; + + // Walk up the parent chain to find the nearest operation with an explicit + // data layout spec. Check ModuleOp explicitly since it does not actually + // implement DataLayoutOpInterface as a trait (it just has the same methods). + Operation *current = op; + while (current) { + // Check for ModuleOp with explicit data layout spec + if (auto mod = llvm::dyn_cast(current)) { + if (mod.getDataLayoutSpec()) + return DataLayout(mod); + } else if (auto dataLayoutOp = + llvm::dyn_cast(current)) { + // Check other DataLayoutOpInterface implementations + if (dataLayoutOp.getDataLayoutSpec()) + return DataLayout(dataLayoutOp); + } + current = current->getParentOp(); + } + + // No explicit data layout found; return default if allowed + if (allowDefault) { + // Check if op itself is a ModuleOp + if (auto mod = llvm::dyn_cast(op)) + return DataLayout(mod); + // Otherwise check parents + if (auto mod = op->getParentOfType()) + return DataLayout(mod); + } + + return std::nullopt; +} + +} // namespace acc +} // namespace mlir diff --git a/mlir/lib/Dialect/OpenACC/Utils/OpenACCUtilsGPU.cpp b/mlir/lib/Dialect/OpenACC/Utils/OpenACCUtilsGPU.cpp new file mode 100644 index 0000000000000..fb1e34f514da9 --- /dev/null +++ b/mlir/lib/Dialect/OpenACC/Utils/OpenACCUtilsGPU.cpp @@ -0,0 +1,47 @@ +//===- OpenACCUtilsGPU.cpp - OpenACC GPU Utilities ------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// +// This file implements utility functions for OpenACC that depend on the GPU +// dialect. +// +//===----------------------------------------------------------------------===// + +#include "mlir/Dialect/OpenACC/OpenACCUtilsGPU.h" +#include "mlir/IR/SymbolTable.h" + +namespace mlir { +namespace acc { + +std::optional getOrCreateGPUModule(ModuleOp mod, bool create, + llvm::StringRef name) { + // Use default name if provided name is empty + llvm::StringRef moduleName = + name.empty() ? llvm::StringRef(kDefaultGPUModuleName) : name; + + // Look for existing GPU module with the specified name + SymbolTable symTab(mod); + if (auto gpuMod = symTab.lookup(moduleName)) + return gpuMod; + + if (!create) + return std::nullopt; + + // Create a new GPU module + auto *ctx = mod.getContext(); + mod->setAttr(gpu::GPUDialect::getContainerModuleAttrName(), + UnitAttr::get(ctx)); + + OpBuilder builder(ctx); + auto gpuMod = gpu::GPUModuleOp::create(builder, mod.getLoc(), moduleName); + Block::iterator insertPt(mod.getBodyRegion().front().end()); + symTab.insert(gpuMod, insertPt); + return gpuMod; +} + +} // namespace acc +} // namespace mlir diff --git a/mlir/unittests/Dialect/OpenACC/CMakeLists.txt b/mlir/unittests/Dialect/OpenACC/CMakeLists.txt index 4359bae00ca21..7303ff581abc4 100644 --- a/mlir/unittests/Dialect/OpenACC/CMakeLists.txt +++ b/mlir/unittests/Dialect/OpenACC/CMakeLists.txt @@ -2,6 +2,8 @@ add_mlir_unittest(MLIROpenACCTests OpenACCOpsTest.cpp OpenACCOpsInterfacesTest.cpp OpenACCTypeInterfacesTest.cpp + OpenACCUtilsCGTest.cpp + OpenACCUtilsGPUTest.cpp OpenACCUtilsTest.cpp OpenACCUtilsTilingTest.cpp OpenACCUtilsLoopTest.cpp @@ -10,6 +12,7 @@ mlir_target_link_libraries(MLIROpenACCTests PRIVATE MLIRIR MLIRAffineDialect + MLIRDLTIDialect MLIRFuncDialect MLIRGPUDialect MLIRMemRefDialect diff --git a/mlir/unittests/Dialect/OpenACC/OpenACCUtilsCGTest.cpp b/mlir/unittests/Dialect/OpenACC/OpenACCUtilsCGTest.cpp new file mode 100644 index 0000000000000..b2d5409f495f5 --- /dev/null +++ b/mlir/unittests/Dialect/OpenACC/OpenACCUtilsCGTest.cpp @@ -0,0 +1,76 @@ +//===- OpenACCUtilsCGTest.cpp - Unit tests for OpenACC CG utilities -------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include "mlir/Dialect/OpenACC/OpenACCUtilsCG.h" +#include "mlir/Dialect/DLTI/DLTI.h" +#include "mlir/Dialect/OpenACC/OpenACC.h" +#include "mlir/IR/BuiltinOps.h" +#include "mlir/IR/MLIRContext.h" +#include "mlir/IR/OwningOpRef.h" +#include "gtest/gtest.h" + +using namespace mlir; +using namespace mlir::acc; + +//===----------------------------------------------------------------------===// +// Test Fixture +//===----------------------------------------------------------------------===// + +class OpenACCUtilsCGTest : public ::testing::Test { +protected: + OpenACCUtilsCGTest() : b(&context), loc(UnknownLoc::get(&context)) { + context.loadDialect(); + } + + MLIRContext context; + OpBuilder b; + Location loc; +}; + +//===----------------------------------------------------------------------===// +// getDataLayout Tests +//===----------------------------------------------------------------------===// + +TEST_F(OpenACCUtilsCGTest, getDataLayoutNoSpecAllowDefault) { + OwningOpRef module = ModuleOp::create(b, loc); + + // With allowDefault=true, should return a default DataLayout + auto dl = getDataLayout(module->getOperation(), /*allowDefault=*/true); + EXPECT_TRUE(dl.has_value()); +} + +TEST_F(OpenACCUtilsCGTest, getDataLayoutNoSpecDisallowDefault) { + OwningOpRef module = ModuleOp::create(b, loc); + + // With allowDefault=false and no spec, should return nullopt + auto dl = getDataLayout(module->getOperation(), /*allowDefault=*/false); + EXPECT_FALSE(dl.has_value()); +} + +TEST_F(OpenACCUtilsCGTest, getDataLayoutNullOp) { + // Null operation should return nullopt + auto dl = getDataLayout(nullptr, /*allowDefault=*/true); + EXPECT_FALSE(dl.has_value()); +} + +TEST_F(OpenACCUtilsCGTest, getDataLayoutWithSpec) { + OwningOpRef module = ModuleOp::create(b, loc); + + // Add a data layout spec to the module + auto indexEntry = DataLayoutEntryAttr::get(IndexType::get(&context), + b.getI32IntegerAttr(32)); + auto spec = DataLayoutSpecAttr::get(&context, {indexEntry}); + (*module)->setAttr(DLTIDialect::kDataLayoutAttrName, spec); + + // With explicit spec, should return DataLayout regardless of allowDefault + auto dl1 = getDataLayout(module->getOperation(), /*allowDefault=*/false); + EXPECT_TRUE(dl1.has_value()); + + auto dl2 = getDataLayout(module->getOperation(), /*allowDefault=*/true); + EXPECT_TRUE(dl2.has_value()); +} diff --git a/mlir/unittests/Dialect/OpenACC/OpenACCUtilsGPUTest.cpp b/mlir/unittests/Dialect/OpenACC/OpenACCUtilsGPUTest.cpp new file mode 100644 index 0000000000000..e9d94edde3a92 --- /dev/null +++ b/mlir/unittests/Dialect/OpenACC/OpenACCUtilsGPUTest.cpp @@ -0,0 +1,89 @@ +//===- OpenACCUtilsGPUTest.cpp - Unit tests for OpenACC GPU utilities ----===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include "mlir/Dialect/OpenACC/OpenACCUtilsGPU.h" +#include "mlir/Dialect/GPU/IR/GPUDialect.h" +#include "mlir/IR/BuiltinOps.h" +#include "mlir/IR/MLIRContext.h" +#include "mlir/IR/OwningOpRef.h" +#include "gtest/gtest.h" + +using namespace mlir; +using namespace mlir::acc; + +//===----------------------------------------------------------------------===// +// Test Fixture +//===----------------------------------------------------------------------===// + +class OpenACCUtilsGPUTest : public ::testing::Test { +protected: + OpenACCUtilsGPUTest() : b(&context), loc(UnknownLoc::get(&context)) { + context.loadDialect(); + } + + MLIRContext context; + OpBuilder b; + Location loc; +}; + +//===----------------------------------------------------------------------===// +// getOrCreateGPUModule Tests +//===----------------------------------------------------------------------===// + +TEST_F(OpenACCUtilsGPUTest, getOrCreateGPUModuleCreatesWhenMissing) { + OwningOpRef module = ModuleOp::create(b, loc); + + // First call should create the GPU module + auto gpuMod = getOrCreateGPUModule(*module, /*create=*/true); + ASSERT_TRUE(gpuMod.has_value()); + EXPECT_EQ(gpuMod->getName(), kDefaultGPUModuleName); + + // Module should now have the container module attribute + EXPECT_TRUE( + (*module)->hasAttr(gpu::GPUDialect::getContainerModuleAttrName())); +} + +TEST_F(OpenACCUtilsGPUTest, getOrCreateGPUModuleReturnsExisting) { + OwningOpRef module = ModuleOp::create(b, loc); + + // Create a GPU module first + auto gpuMod1 = getOrCreateGPUModule(*module, /*create=*/true); + ASSERT_TRUE(gpuMod1.has_value()); + + // Second call should return the same module + auto gpuMod2 = getOrCreateGPUModule(*module, /*create=*/true); + ASSERT_TRUE(gpuMod2.has_value()); + EXPECT_EQ(gpuMod1->getOperation(), gpuMod2->getOperation()); +} + +TEST_F(OpenACCUtilsGPUTest, getOrCreateGPUModuleNoCreateReturnsNullopt) { + OwningOpRef module = ModuleOp::create(b, loc); + + // With create=false and no existing GPU module, should return nullopt + auto gpuMod = getOrCreateGPUModule(*module, /*create=*/false); + EXPECT_FALSE(gpuMod.has_value()); +} + +TEST_F(OpenACCUtilsGPUTest, getOrCreateGPUModuleCustomName) { + OwningOpRef module = ModuleOp::create(b, loc); + + // Create with custom name + auto gpuMod = + getOrCreateGPUModule(*module, /*create=*/true, "custom_gpu_module"); + ASSERT_TRUE(gpuMod.has_value()); + EXPECT_EQ(gpuMod->getName(), "custom_gpu_module"); +} + +TEST_F(OpenACCUtilsGPUTest, getOrCreateGPUModuleEmptyNameUsesDefault) { + OwningOpRef module = ModuleOp::create(b, loc); + + // Empty name should use default + auto gpuMod = getOrCreateGPUModule(*module, /*create=*/true, ""); + ASSERT_TRUE(gpuMod.has_value()); + EXPECT_EQ(gpuMod->getName(), kDefaultGPUModuleName); +} From 1d18539a7b60c05a587ac439b1d9095d01ac691e Mon Sep 17 00:00:00 2001 From: Jay Foad Date: Wed, 11 Feb 2026 19:14:54 +0000 Subject: [PATCH 002/705] [AMDGPU] Treat F64 TRANS instructions as VALU for S_DELAY_ALU insertion (#180940) This matches the hardware behavior. --- .../Target/AMDGPU/AMDGPUInsertDelayAlu.cpp | 4 +++- .../CodeGen/AMDGPU/GlobalISel/fdiv.f64.ll | 12 +++++------ .../GlobalISel/llvm.amdgcn.rsq.clamp.ll | 12 ++++------- llvm/test/CodeGen/AMDGPU/frem.ll | 20 +++++++++---------- llvm/test/CodeGen/AMDGPU/insert-delay-alu.mir | 17 ++++++++++++++++ 5 files changed, 40 insertions(+), 25 deletions(-) diff --git a/llvm/lib/Target/AMDGPU/AMDGPUInsertDelayAlu.cpp b/llvm/lib/Target/AMDGPU/AMDGPUInsertDelayAlu.cpp index 7c84edf2a60bc..d1b9fb4115f66 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPUInsertDelayAlu.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPUInsertDelayAlu.cpp @@ -68,7 +68,9 @@ class AMDGPUInsertDelayAlu { // Get the delay type for a MachineInstr. DelayType getDelayType(const MachineInstr &MI) { - if (SIInstrInfo::isTRANS(MI)) + // Non-F64 TRANS instructions use a separate delay type. + if (SIInstrInfo::isTRANS(MI) && + !AMDGPU::isDPMACCInstruction(MI.getOpcode())) return TRANS; // WMMA XDL ops are treated the same as TRANS. if (ST->hasGFX1250Insts() && SII->isXDLWMMA(MI)) diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/fdiv.f64.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/fdiv.f64.ll index 31fc1ee538957..3f6a6ac44f543 100644 --- a/llvm/test/CodeGen/AMDGPU/GlobalISel/fdiv.f64.ll +++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/fdiv.f64.ll @@ -832,7 +832,7 @@ define <2 x double> @v_fdiv_v2f64(<2 x double> %a, <2 x double> %b) { ; GFX11-NEXT: v_div_scale_f64 v[8:9], null, v[4:5], v[4:5], v[0:1] ; GFX11-NEXT: v_div_scale_f64 v[10:11], null, v[6:7], v[6:7], v[2:3] ; GFX11-NEXT: v_div_scale_f64 v[20:21], vcc_lo, v[0:1], v[4:5], v[0:1] -; GFX11-NEXT: s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2) +; GFX11-NEXT: s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3) ; GFX11-NEXT: v_rcp_f64_e32 v[12:13], v[8:9] ; GFX11-NEXT: v_rcp_f64_e32 v[14:15], v[10:11] ; GFX11-NEXT: s_waitcnt_depctr depctr_va_vdst(0) @@ -1065,7 +1065,7 @@ define <2 x double> @v_fdiv_v2f64_ulp25(<2 x double> %a, <2 x double> %b) { ; GFX11-NEXT: v_div_scale_f64 v[8:9], null, v[4:5], v[4:5], v[0:1] ; GFX11-NEXT: v_div_scale_f64 v[10:11], null, v[6:7], v[6:7], v[2:3] ; GFX11-NEXT: v_div_scale_f64 v[20:21], vcc_lo, v[0:1], v[4:5], v[0:1] -; GFX11-NEXT: s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2) +; GFX11-NEXT: s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3) ; GFX11-NEXT: v_rcp_f64_e32 v[12:13], v[8:9] ; GFX11-NEXT: v_rcp_f64_e32 v[14:15], v[10:11] ; GFX11-NEXT: s_waitcnt_depctr depctr_va_vdst(0) @@ -1225,7 +1225,7 @@ define <2 x double> @v_rcp_v2f64(<2 x double> %x) { ; GFX11-NEXT: v_div_scale_f64 v[4:5], null, v[0:1], v[0:1], 1.0 ; GFX11-NEXT: v_div_scale_f64 v[6:7], null, v[2:3], v[2:3], 1.0 ; GFX11-NEXT: v_div_scale_f64 v[16:17], vcc_lo, 1.0, v[0:1], 1.0 -; GFX11-NEXT: s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2) +; GFX11-NEXT: s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3) ; GFX11-NEXT: v_rcp_f64_e32 v[8:9], v[4:5] ; GFX11-NEXT: v_rcp_f64_e32 v[10:11], v[6:7] ; GFX11-NEXT: s_waitcnt_depctr depctr_va_vdst(0) @@ -1385,7 +1385,7 @@ define <2 x double> @v_rcp_v2f64_arcp(<2 x double> %x) { ; GFX11-NEXT: v_div_scale_f64 v[4:5], null, v[0:1], v[0:1], 1.0 ; GFX11-NEXT: v_div_scale_f64 v[6:7], null, v[2:3], v[2:3], 1.0 ; GFX11-NEXT: v_div_scale_f64 v[16:17], vcc_lo, 1.0, v[0:1], 1.0 -; GFX11-NEXT: s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2) +; GFX11-NEXT: s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3) ; GFX11-NEXT: v_rcp_f64_e32 v[8:9], v[4:5] ; GFX11-NEXT: v_rcp_f64_e32 v[10:11], v[6:7] ; GFX11-NEXT: s_waitcnt_depctr depctr_va_vdst(0) @@ -1612,7 +1612,7 @@ define <2 x double> @v_rcp_v2f64_ulp25(<2 x double> %x) { ; GFX11-NEXT: v_div_scale_f64 v[4:5], null, v[0:1], v[0:1], 1.0 ; GFX11-NEXT: v_div_scale_f64 v[6:7], null, v[2:3], v[2:3], 1.0 ; GFX11-NEXT: v_div_scale_f64 v[16:17], vcc_lo, 1.0, v[0:1], 1.0 -; GFX11-NEXT: s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2) +; GFX11-NEXT: s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3) ; GFX11-NEXT: v_rcp_f64_e32 v[8:9], v[4:5] ; GFX11-NEXT: v_rcp_f64_e32 v[10:11], v[6:7] ; GFX11-NEXT: s_waitcnt_depctr depctr_va_vdst(0) @@ -1845,7 +1845,7 @@ define <2 x double> @v_fdiv_v2f64_arcp_ulp25(<2 x double> %a, <2 x double> %b) { ; GFX11-NEXT: v_div_scale_f64 v[8:9], null, v[4:5], v[4:5], v[0:1] ; GFX11-NEXT: v_div_scale_f64 v[10:11], null, v[6:7], v[6:7], v[2:3] ; GFX11-NEXT: v_div_scale_f64 v[20:21], vcc_lo, v[0:1], v[4:5], v[0:1] -; GFX11-NEXT: s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_2) +; GFX11-NEXT: s_delay_alu instid0(VALU_DEP_3) | instskip(NEXT) | instid1(VALU_DEP_3) ; GFX11-NEXT: v_rcp_f64_e32 v[12:13], v[8:9] ; GFX11-NEXT: v_rcp_f64_e32 v[14:15], v[10:11] ; GFX11-NEXT: s_waitcnt_depctr depctr_va_vdst(0) diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.rsq.clamp.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.rsq.clamp.ll index 26cdbb11d27d6..4a6e24b700663 100644 --- a/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.rsq.clamp.ll +++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.rsq.clamp.ll @@ -95,11 +95,10 @@ define double @v_rsq_clamp_f64(double %src) #0 { ; GFX12-NEXT: v_rsq_f64_e32 v[0:1], v[0:1] ; GFX12-NEXT: v_mov_b32_e32 v2, -1 ; GFX12-NEXT: v_mov_b32_e32 v3, 0x7fefffff -; GFX12-NEXT: s_delay_alu instid0(TRANS32_DEP_1) | instid1(VALU_DEP_1) +; GFX12-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1) ; GFX12-NEXT: v_min_num_f64_e32 v[0:1], v[0:1], v[2:3] ; GFX12-NEXT: v_mov_b32_e32 v2, -1 ; GFX12-NEXT: v_mov_b32_e32 v3, 0xffefffff -; GFX12-NEXT: s_delay_alu instid0(VALU_DEP_1) ; GFX12-NEXT: v_max_num_f64_e32 v[0:1], v[0:1], v[2:3] ; GFX12-NEXT: s_setpc_b64 s[30:31] %rsq_clamp = call double @llvm.amdgcn.rsq.clamp.f64(double %src) @@ -135,11 +134,10 @@ define double @v_rsq_clamp_fabs_f64(double %src) #0 { ; GFX12-NEXT: v_rsq_f64_e64 v[0:1], |v[0:1]| ; GFX12-NEXT: v_mov_b32_e32 v2, -1 ; GFX12-NEXT: v_mov_b32_e32 v3, 0x7fefffff -; GFX12-NEXT: s_delay_alu instid0(TRANS32_DEP_1) | instid1(VALU_DEP_1) +; GFX12-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1) ; GFX12-NEXT: v_min_num_f64_e32 v[0:1], v[0:1], v[2:3] ; GFX12-NEXT: v_mov_b32_e32 v2, -1 ; GFX12-NEXT: v_mov_b32_e32 v3, 0xffefffff -; GFX12-NEXT: s_delay_alu instid0(VALU_DEP_1) ; GFX12-NEXT: v_max_num_f64_e32 v[0:1], v[0:1], v[2:3] ; GFX12-NEXT: s_setpc_b64 s[30:31] %fabs.src = call double @llvm.fabs.f64(double %src) @@ -208,11 +206,10 @@ define double @v_rsq_clamp_undef_f64() #0 { ; GFX12-NEXT: v_rsq_f64_e32 v[0:1], s[0:1] ; GFX12-NEXT: v_mov_b32_e32 v2, -1 ; GFX12-NEXT: v_mov_b32_e32 v3, 0x7fefffff -; GFX12-NEXT: s_delay_alu instid0(TRANS32_DEP_1) | instid1(VALU_DEP_1) +; GFX12-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1) ; GFX12-NEXT: v_min_num_f64_e32 v[0:1], v[0:1], v[2:3] ; GFX12-NEXT: v_mov_b32_e32 v2, -1 ; GFX12-NEXT: v_mov_b32_e32 v3, 0xffefffff -; GFX12-NEXT: s_delay_alu instid0(VALU_DEP_1) ; GFX12-NEXT: v_max_num_f64_e32 v[0:1], v[0:1], v[2:3] ; GFX12-NEXT: s_setpc_b64 s[30:31] %rsq_clamp = call double @llvm.amdgcn.rsq.clamp.f64(double poison) @@ -279,11 +276,10 @@ define double @v_rsq_clamp_f64_non_ieee(double %src) #2 { ; GFX12-NEXT: v_rsq_f64_e32 v[0:1], v[0:1] ; GFX12-NEXT: v_mov_b32_e32 v2, -1 ; GFX12-NEXT: v_mov_b32_e32 v3, 0x7fefffff -; GFX12-NEXT: s_delay_alu instid0(TRANS32_DEP_1) | instid1(VALU_DEP_1) +; GFX12-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(VALU_DEP_1) ; GFX12-NEXT: v_min_num_f64_e32 v[0:1], v[0:1], v[2:3] ; GFX12-NEXT: v_mov_b32_e32 v2, -1 ; GFX12-NEXT: v_mov_b32_e32 v3, 0xffefffff -; GFX12-NEXT: s_delay_alu instid0(VALU_DEP_1) ; GFX12-NEXT: v_max_num_f64_e32 v[0:1], v[0:1], v[2:3] ; GFX12-NEXT: s_setpc_b64 s[30:31] %rsq_clamp = call double @llvm.amdgcn.rsq.clamp.f64(double %src) diff --git a/llvm/test/CodeGen/AMDGPU/frem.ll b/llvm/test/CodeGen/AMDGPU/frem.ll index 5d243189d516a..bf153a88982e0 100644 --- a/llvm/test/CodeGen/AMDGPU/frem.ll +++ b/llvm/test/CodeGen/AMDGPU/frem.ll @@ -4224,7 +4224,7 @@ define amdgpu_kernel void @frem_f64(ptr addrspace(1) %out, ptr addrspace(1) %in1 ; GFX1150-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) ; GFX1150-NEXT: v_div_scale_f64 v[8:9], null, v[4:5], v[4:5], 1.0 ; GFX1150-NEXT: v_rcp_f64_e32 v[10:11], v[8:9] -; GFX1150-NEXT: s_delay_alu instid0(TRANS32_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) +; GFX1150-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) ; GFX1150-NEXT: v_fma_f64 v[14:15], -v[8:9], v[10:11], 1.0 ; GFX1150-NEXT: v_fma_f64 v[10:11], v[10:11], v[14:15], v[10:11] ; GFX1150-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) @@ -4330,7 +4330,7 @@ define amdgpu_kernel void @frem_f64(ptr addrspace(1) %out, ptr addrspace(1) %in1 ; GFX1200-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) ; GFX1200-NEXT: v_div_scale_f64 v[8:9], null, v[4:5], v[4:5], 1.0 ; GFX1200-NEXT: v_rcp_f64_e32 v[10:11], v[8:9] -; GFX1200-NEXT: s_delay_alu instid0(TRANS32_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) +; GFX1200-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) ; GFX1200-NEXT: v_fma_f64 v[14:15], -v[8:9], v[10:11], 1.0 ; GFX1200-NEXT: v_fma_f64 v[10:11], v[10:11], v[14:15], v[10:11] ; GFX1200-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) @@ -4622,7 +4622,7 @@ define amdgpu_kernel void @fast_frem_f64(ptr addrspace(1) %out, ptr addrspace(1) ; GFX1150-NEXT: global_load_b64 v[2:3], v12, s[4:5] ; GFX1150-NEXT: s_waitcnt vmcnt(0) ; GFX1150-NEXT: v_div_scale_f64 v[4:5], null, v[2:3], v[2:3], v[0:1] -; GFX1150-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(TRANS32_DEP_1) +; GFX1150-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) ; GFX1150-NEXT: v_rcp_f64_e32 v[6:7], v[4:5] ; GFX1150-NEXT: v_fma_f64 v[8:9], -v[4:5], v[6:7], 1.0 ; GFX1150-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) @@ -4655,7 +4655,7 @@ define amdgpu_kernel void @fast_frem_f64(ptr addrspace(1) %out, ptr addrspace(1) ; GFX1200-NEXT: global_load_b64 v[2:3], v12, s[4:5] ; GFX1200-NEXT: s_wait_loadcnt 0x0 ; GFX1200-NEXT: v_div_scale_f64 v[4:5], null, v[2:3], v[2:3], v[0:1] -; GFX1200-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(TRANS32_DEP_1) +; GFX1200-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) ; GFX1200-NEXT: v_rcp_f64_e32 v[6:7], v[4:5] ; GFX1200-NEXT: v_fma_f64 v[8:9], -v[4:5], v[6:7], 1.0 ; GFX1200-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) @@ -4900,7 +4900,7 @@ define amdgpu_kernel void @unsafe_frem_f64(ptr addrspace(1) %out, ptr addrspace( ; GFX1150-NEXT: global_load_b64 v[2:3], v12, s[4:5] ; GFX1150-NEXT: s_waitcnt vmcnt(0) ; GFX1150-NEXT: v_div_scale_f64 v[4:5], null, v[2:3], v[2:3], v[0:1] -; GFX1150-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(TRANS32_DEP_1) +; GFX1150-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) ; GFX1150-NEXT: v_rcp_f64_e32 v[6:7], v[4:5] ; GFX1150-NEXT: v_fma_f64 v[8:9], -v[4:5], v[6:7], 1.0 ; GFX1150-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) @@ -4933,7 +4933,7 @@ define amdgpu_kernel void @unsafe_frem_f64(ptr addrspace(1) %out, ptr addrspace( ; GFX1200-NEXT: global_load_b64 v[2:3], v12, s[4:5] ; GFX1200-NEXT: s_wait_loadcnt 0x0 ; GFX1200-NEXT: v_div_scale_f64 v[4:5], null, v[2:3], v[2:3], v[0:1] -; GFX1200-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(TRANS32_DEP_1) +; GFX1200-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) ; GFX1200-NEXT: v_rcp_f64_e32 v[6:7], v[4:5] ; GFX1200-NEXT: v_fma_f64 v[8:9], -v[4:5], v[6:7], 1.0 ; GFX1200-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) @@ -17172,7 +17172,7 @@ define amdgpu_kernel void @frem_v2f64(ptr addrspace(1) %out, ptr addrspace(1) %i ; GFX1150-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) ; GFX1150-NEXT: v_div_scale_f64 v[12:13], null, v[8:9], v[8:9], 1.0 ; GFX1150-NEXT: v_rcp_f64_e32 v[14:15], v[12:13] -; GFX1150-NEXT: s_delay_alu instid0(TRANS32_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) +; GFX1150-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) ; GFX1150-NEXT: v_fma_f64 v[18:19], -v[12:13], v[14:15], 1.0 ; GFX1150-NEXT: v_fma_f64 v[14:15], v[14:15], v[18:19], v[14:15] ; GFX1150-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) @@ -17259,7 +17259,7 @@ define amdgpu_kernel void @frem_v2f64(ptr addrspace(1) %out, ptr addrspace(1) %i ; GFX1150-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) ; GFX1150-NEXT: v_div_scale_f64 v[14:15], null, v[10:11], v[10:11], 1.0 ; GFX1150-NEXT: v_rcp_f64_e32 v[16:17], v[14:15] -; GFX1150-NEXT: s_delay_alu instid0(TRANS32_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) +; GFX1150-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) ; GFX1150-NEXT: v_fma_f64 v[20:21], -v[14:15], v[16:17], 1.0 ; GFX1150-NEXT: v_fma_f64 v[16:17], v[16:17], v[20:21], v[16:17] ; GFX1150-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) @@ -17371,7 +17371,7 @@ define amdgpu_kernel void @frem_v2f64(ptr addrspace(1) %out, ptr addrspace(1) %i ; GFX1200-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) ; GFX1200-NEXT: v_div_scale_f64 v[12:13], null, v[8:9], v[8:9], 1.0 ; GFX1200-NEXT: v_rcp_f64_e32 v[14:15], v[12:13] -; GFX1200-NEXT: s_delay_alu instid0(TRANS32_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) +; GFX1200-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) ; GFX1200-NEXT: v_fma_f64 v[18:19], -v[12:13], v[14:15], 1.0 ; GFX1200-NEXT: v_fma_f64 v[14:15], v[14:15], v[18:19], v[14:15] ; GFX1200-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) @@ -17462,7 +17462,7 @@ define amdgpu_kernel void @frem_v2f64(ptr addrspace(1) %out, ptr addrspace(1) %i ; GFX1200-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) ; GFX1200-NEXT: v_div_scale_f64 v[14:15], null, v[10:11], v[10:11], 1.0 ; GFX1200-NEXT: v_rcp_f64_e32 v[16:17], v[14:15] -; GFX1200-NEXT: s_delay_alu instid0(TRANS32_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) +; GFX1200-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) ; GFX1200-NEXT: v_fma_f64 v[20:21], -v[14:15], v[16:17], 1.0 ; GFX1200-NEXT: v_fma_f64 v[16:17], v[16:17], v[20:21], v[16:17] ; GFX1200-NEXT: s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1) diff --git a/llvm/test/CodeGen/AMDGPU/insert-delay-alu.mir b/llvm/test/CodeGen/AMDGPU/insert-delay-alu.mir index c287fb3614496..3addab3508140 100644 --- a/llvm/test/CodeGen/AMDGPU/insert-delay-alu.mir +++ b/llvm/test/CodeGen/AMDGPU/insert-delay-alu.mir @@ -379,6 +379,15 @@ ; CHECK-NEXT: v_mul_f32_e64 v0, v0, v0 ret void; } + + define void @trans64() { + ; CHECK-LABEL: trans64: + ; CHECK: ; %bb.0: + ; CHECK-NEXT: v_rcp_f64_e32 v[0:1], v[0:1] + ; CHECK-NEXT: s_delay_alu instid0(VALU_DEP_1) + ; CHECK-NEXT: v_add_f64 v[0:1], v[0:1], v[0:1] + ret void; + } ... --- @@ -762,3 +771,11 @@ body: | $vgpr0 = V_MUL_F32_e64 0, $vgpr0, 0, $vgpr0, 0, 0, implicit $mode, implicit $exec ... +# Check that F64 TRANS instructions are treated as VALU. +--- +name: trans64 +body: | + bb.0: + $vgpr0_vgpr1 = V_RCP_F64_e32 $vgpr0_vgpr1, implicit $exec, implicit $mode + $vgpr0_vgpr1 = V_ADD_F64_e64 0, $vgpr0_vgpr1, 0, $vgpr0_vgpr1, 0, 0, implicit $exec, implicit $mode +... From a64f541d9eb1720b5224de38a46cd40bc18a9eba Mon Sep 17 00:00:00 2001 From: Louis Dionne Date: Wed, 11 Feb 2026 14:15:20 -0500 Subject: [PATCH 003/705] [libc++] Add scripts to run benchmarks and submit to LNT on a regular basis (#180849) Also add a libc++ schema that can be used to create a test suite capable of storing libc++ benchmarking data on any LNT instance. --- libcxx/utils/ci/benchmark-for-lnt.py | 133 -------------------- libcxx/utils/ci/lnt/README.md | 31 +++++ libcxx/utils/ci/lnt/commit-watch | 132 +++++++++++++++++++ libcxx/utils/ci/lnt/run-benchmarks | 182 +++++++++++++++++++++++++++ libcxx/utils/ci/lnt/schema.yaml | 44 +++++++ 5 files changed, 389 insertions(+), 133 deletions(-) delete mode 100755 libcxx/utils/ci/benchmark-for-lnt.py create mode 100644 libcxx/utils/ci/lnt/README.md create mode 100755 libcxx/utils/ci/lnt/commit-watch create mode 100755 libcxx/utils/ci/lnt/run-benchmarks create mode 100644 libcxx/utils/ci/lnt/schema.yaml diff --git a/libcxx/utils/ci/benchmark-for-lnt.py b/libcxx/utils/ci/benchmark-for-lnt.py deleted file mode 100755 index a139c58bb366f..0000000000000 --- a/libcxx/utils/ci/benchmark-for-lnt.py +++ /dev/null @@ -1,133 +0,0 @@ -#!/usr/bin/env python -# ===----------------------------------------------------------------------===## -# -# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -# See https://llvm.org/LICENSE.txt for license information. -# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -# -# ===----------------------------------------------------------------------===## - -import argparse -import os -import pathlib -import subprocess -import sys -import tempfile - -def step(message: str) -> None: - print(message, file=sys.stderr) - -def directory_path(string): - if os.path.isdir(string): - return pathlib.Path(string) - else: - raise NotADirectoryError(string) - -def main(argv): - parser = argparse.ArgumentParser( - prog='benchmark-for-lnt', - description='Benchmark libc++ at the given commit for submitting to LNT.') - parser.add_argument('-o', '--output', type=argparse.FileType('w'), default='-', - help='Path to the file where the resulting LNT report containing benchmark results is written. ' - 'By default, stdout.') - parser.add_argument('--benchmark-commit', type=str, required=True, - help='The SHA representing the version of the library to benchmark.') - parser.add_argument('--test-suite-commit', type=str, required=True, - help='The SHA representing the version of the test suite to use for benchmarking.') - parser.add_argument('--machine', type=str, required=True, - help='The name of the machine for reporting LNT results.') - parser.add_argument('--spec-dir', type=pathlib.Path, required=False, - help='Optional path to a SPEC installation to use for benchmarking.') - parser.add_argument('--git-repo', type=directory_path, default=os.getcwd(), - help='Optional path to the Git repository to use. By default, the current working directory is used.') - parser.add_argument('--dry-run', action='store_true', - help='Only print what would be executed.') - parser.add_argument('-v', '--verbose', action='store_true', - help='Print the output of all subcommands.') - args = parser.parse_args(argv) - - def run(command, *posargs, enforce_success=True, **kwargs): - command = [str(c) for c in command] - if args.dry_run: - print(f'$ {" ".join(command)}') - else: - # If we're running with verbose, print everything but redirect output to stderr since - # we already output the json to stdout in some cases. Otherwise, hush everything. - if args.verbose: - if 'stdout' not in kwargs: - kwargs.update({'stdout': sys.stderr}) - else: - if 'stdout' not in kwargs: - kwargs.update({'stdout': subprocess.DEVNULL}) - if 'stderr' not in kwargs: - kwargs.update({'stderr': subprocess.DEVNULL}) - if enforce_success: - subprocess.check_call(command, *posargs, **kwargs) - else: - subprocess.call(command, *posargs, **kwargs) - - with tempfile.TemporaryDirectory() as build_dir: - build_dir = pathlib.Path(build_dir) - - step(f'Building libc++ at commit {args.benchmark_commit}') - run([args.git_repo / 'libcxx/utils/build-at-commit', - '--git-repo', args.git_repo, - '--install-dir', build_dir / 'install', - '--commit', args.benchmark_commit, - '--', '-DCMAKE_BUILD_TYPE=RelWithDebInfo']) - - if args.spec_dir is not None: - step(f'Running SPEC benchmarks from {args.test_suite_commit} against libc++ {args.benchmark_commit}') - run([args.git_repo / 'libcxx/utils/test-at-commit', - '--git-repo', args.git_repo, - '--build-dir', build_dir / 'spec', - '--test-suite-commit', args.test_suite_commit, - '--libcxx-installation', build_dir / 'install', - '--', - '-j1', '--time-tests', - '--param', 'optimization=speed', - '--param', 'std=c++17', - '--param', f'spec_dir={args.spec_dir}', - build_dir / 'spec/libcxx/test', - '--filter', 'benchmarks/spec.gen.py'], - enforce_success=False) - - # TODO: For now, we run only a subset of the benchmarks because running the whole test suite is too slow. - # Run the whole test suite once https://github.com/llvm/llvm-project/issues/173032 is resolved. - step(f'Running microbenchmarks from {args.test_suite_commit} against libc++ {args.benchmark_commit}') - run([args.git_repo / 'libcxx/utils/test-at-commit', - '--git-repo', args.git_repo, - '--build-dir', build_dir / 'micro', - '--test-suite-commit', args.test_suite_commit, - '--libcxx-installation', build_dir / 'install', - '--', - '-j1', '--time-tests', - '--param', 'optimization=speed', - '--param', 'std=c++26', - build_dir / 'micro/libcxx/test', - '--filter', 'benchmarks/(algorithms|containers|iterators|locale|memory|streams|numeric|utility)'], - enforce_success=False) - - step('Installing LNT') - run(['python', '-m', 'venv', build_dir / '.venv']) - run([build_dir / '.venv/bin/pip', 'install', 'llvm-lnt']) - - step('Consolidating benchmark results and creating JSON report') - if args.spec_dir is not None: - with open(build_dir / 'benchmarks.lnt', 'w') as f: - run([args.git_repo / 'libcxx/utils/consolidate-benchmarks', build_dir / 'spec'], stdout=f) - with open(build_dir / 'benchmarks.lnt', 'a') as f: - run([args.git_repo / 'libcxx/utils/consolidate-benchmarks', build_dir / 'micro'], stdout=f) - order = len(subprocess.check_output(['git', '-C', args.git_repo, 'rev-list', args.benchmark_commit]).splitlines()) - commit_info = subprocess.check_output(['git', '-C', args.git_repo, 'show', args.benchmark_commit, '--no-patch']).decode() - run([build_dir / '.venv/bin/lnt', 'importreport', '--order', str(order), '--machine', args.machine, - '--run-info', f'commit_info={commit_info}', - build_dir / 'benchmarks.lnt', build_dir / 'benchmarks.json']) - - if not args.dry_run: - with open(build_dir / 'benchmarks.json', 'r') as f: - args.output.write(f.read()) - - -if __name__ == '__main__': - main(sys.argv[1:]) diff --git a/libcxx/utils/ci/lnt/README.md b/libcxx/utils/ci/lnt/README.md new file mode 100644 index 0000000000000..2f2f42e89d6db --- /dev/null +++ b/libcxx/utils/ci/lnt/README.md @@ -0,0 +1,31 @@ +This directory contains utilities for continuous benchmarking of libc++ with LNT. +This can be done locally using a local instance, or using a public instance like http://lnt.llvm.org. + +Example for running locally: + +``` +# Create an instance and run a server +lnt create my-instance +echo "api_auth_token = 'example_token'" >> my-instance/lnt.cfg +lnt runserver my-instance + +# In another terminal, create the libcxx test suite on the locally-running server +cat < lnt-admin-config.yaml +lnt_url: "http://localhost:8000" +database: default +auth_token: example_token +EOF +lnt admin --config lnt-admin-config.yaml --testsuite libcxx test-suite add libcxx/utils/ci/lnt/schema.yaml + +# Then, watch for libc++ commits and submit benchmark results to the locally-running instance +libcxx/utils/ci/lnt/commit-watch --lnt-url http://localhost:8000 --test-suite libcxx --machine my-laptop | \ + while read commit; do \ + libcxx/utils/ci/lnt/run-benchmarks \ + --test-suite-commit abcdef09 \ + --lnt-url http://localhost:8000 \ + --machine my-laptop \ + --test-suite libcxx \ + --compiler clang++ \ + --benchmark-commit ${commit} \ + done +``` diff --git a/libcxx/utils/ci/lnt/commit-watch b/libcxx/utils/ci/lnt/commit-watch new file mode 100755 index 0000000000000..19aa76cb191fe --- /dev/null +++ b/libcxx/utils/ci/lnt/commit-watch @@ -0,0 +1,132 @@ +#!/usr/bin/env python +# ===----------------------------------------------------------------------===## +# +# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +# See https://llvm.org/LICENSE.txt for license information. +# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +# +# ===----------------------------------------------------------------------===## + +from typing import List, Optional, Set +import argparse +import json +import logging +import os +import pathlib +import subprocess +import sys +import time + + +def directory_path(string): + if os.path.isdir(string): + return pathlib.Path(string) + else: + raise NotADirectoryError(string) + +def api(lnt_url: str, test_suite: str, endpoint: str): + url = f'{lnt_url}/api/db_default/v4/{test_suite}{endpoint}' + logging.debug(f'Querying {url}') + result = json.loads(subprocess.check_output(['curl', '-sS', url]).decode()) + return result + +def get_benchmarked_commits(lnt_url: str, test_suite: str, machine: str) -> Set[str]: + """ + Return the set of commits that have already been benchmarked on the given LNT + instance, test suite and machine. + """ + result = api(lnt_url, test_suite, f'/machines/{machine}') + commits = set() + if 'runs' not in result: # there is no such machine + return set() + for run in result['runs']: + if 'git_sha' not in run: + raise ValueError(f'Found run without a git_sha field: are you using the right LNT test suite? {run}') + commits.add(run['git_sha']) + return commits + +def git_rev_list(git_repo: str, paths: List[str] = []) -> List[str]: + """ + Return the list of all revisions in the given Git repository. Older commits are earlier in the list. + """ + cli = ['git', '-C', git_repo, 'rev-list', 'origin', '--', *paths] + rev_list = subprocess.check_output(cli).decode().strip().splitlines() + return list(reversed(rev_list)) + +def git_sort_revlist(git_repo: str, commits: List[str]) -> List[str]: + """ + Return the list of commits sorted by their chronological order (from oldest to newest) in the + provided Git repository. Items earlier in the list are older than items later in the list. + """ + revlist_cmd = ['git', '-C', git_repo, 'rev-list', '--no-walk'] + list(commits) + revlist = subprocess.check_output(revlist_cmd, text=True).strip().splitlines() + return list(reversed(revlist)) + +def get_all_libcxx_commits(git_repo: str) -> Set[str]: + """ + Return the set of commits available to benchmark for libc++: this is the list of + commits that touch code in libc++ that we care to benchmark. + """ + logging.debug(f'Fetching {git_repo}') + subprocess.check_call(['git', '-C', git_repo, 'fetch', '--quiet', 'origin']) + return set(git_rev_list(git_repo, paths=['libcxx/include', 'libcxx/src'])) + + +def main(argv): + parser = argparse.ArgumentParser( + prog='commit-watch', + description='Watch for libc++ commits to run benchmarks on and print them to standard output.', + formatter_class=argparse.ArgumentDefaultsHelpFormatter) + parser.add_argument('--lnt-url', type=str, required=True, + help='The URL of the LNT instance to use as the source of truth for finding already-benchmarked commits.') + parser.add_argument('--test-suite', type=str, required=True, + help='The name of the test suite that we are producing results for on the LNT instance.') + parser.add_argument('--machine', type=str, required=True, + help='The name of the machine that we are producing results for on the LNT instance.') + parser.add_argument('-d', '--delay', type=int, required=False, default=60, + help='The delay (in seconds) to wait before successive polls.') + parser.add_argument('--git-repo', type=directory_path, default=os.getcwd(), + help='Optional path to the Git repository to use. By default, the current working directory is used.') + parser.add_argument('-v', '--verbose', action='count', default=0, + help='Verbosity level: passing the option multiple times increases the level.') + args = parser.parse_args(argv) + + VERY_VERBOSE = logging.DEBUG - 1 + if args.verbose == 1: + logging.basicConfig(level=logging.INFO) + elif args.verbose == 2: + logging.basicConfig(level=logging.DEBUG) + elif args.verbose > 2: + logging.basicConfig(level=VERY_VERBOSE) + + last_commit: Optional[str] = None # last commit printed + while True: + logging.info(f'Fetching all libc++ commits in {args.git_repo}') + all_commits = get_all_libcxx_commits(args.git_repo) + logging.log(VERY_VERBOSE, f'Found all libc++ commits: {all_commits}') + + logging.info(f'Getting benchmarked commits for {args.machine} from {args.lnt_url}') + benchmarked_commits = get_benchmarked_commits(args.lnt_url, args.test_suite, args.machine) + logging.log(VERY_VERBOSE, f'Found benchmarked commits: {benchmarked_commits}') + + # Print the next commit that should be benchmarked, but only when it changes over the one + # that we just printed. This ensures that we don't repeatedly print that the same benchmark + # should be benchmarked, which would otherwise happen when we are running the benchmarks but + # have not submitted to LNT yet. + commits_to_benchmark = all_commits - benchmarked_commits + if not commits_to_benchmark: + logging.info('No more commits to benchmark') + else: + commits_to_benchmark = git_sort_revlist(args.git_repo, list(commits_to_benchmark)) + most_recent = commits_to_benchmark[-1] + if most_recent != last_commit: + logging.info(f'Selected {most_recent}') + print(most_recent, flush=True) + last_commit = most_recent + + logging.debug(f'Waiting {args.delay}s before next poll') + time.sleep(args.delay) + + +if __name__ == '__main__': + main(sys.argv[1:]) diff --git a/libcxx/utils/ci/lnt/run-benchmarks b/libcxx/utils/ci/lnt/run-benchmarks new file mode 100755 index 0000000000000..65b07e97b8982 --- /dev/null +++ b/libcxx/utils/ci/lnt/run-benchmarks @@ -0,0 +1,182 @@ +#!/usr/bin/env python +# ===----------------------------------------------------------------------===## +# +# Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +# See https://llvm.org/LICENSE.txt for license information. +# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +# +# ===----------------------------------------------------------------------===## + +import argparse +import json +import logging +import os +import pathlib +import platform +import subprocess +import sys +import tempfile + + +def directory_path(string): + if os.path.isdir(string): + return pathlib.Path(string) + else: + raise NotADirectoryError(string) + +def gather_machine_information(args): + """ + Gather the machine information to upload to LNT as part of the submission. + """ + info = {} + if platform.system() == 'Darwin': + profiler_info = json.loads(subprocess.check_output(['system_profiler', 'SPHardwareDataType', 'SPSoftwareDataType', '-json']).decode()) + info['hardware'] = profiler_info['SPHardwareDataType'][0]['chip_type'] + info['os'] = profiler_info['SPSoftwareDataType'][0]['os_version'] + info['sdk'] = subprocess.check_output(['xcrun', '--show-sdk-version']).decode().strip() + + info['compiler'] = subprocess.check_output([args.compiler, '--version']).decode().strip().splitlines()[0] + info['test_suite_commit'] = subprocess.check_output(['git', '-C', args.git_repo, 'rev-parse', args.test_suite_commit]).decode().strip() + + return info + +def gather_run_information(args): + """ + Gather the run information to upload to LNT as part of the submission. + """ + info = {} + info['commit_info'] = subprocess.check_output(['git', '-C', args.git_repo, 'show', args.benchmark_commit, '--no-patch']).decode() + info['git_sha'] = subprocess.check_output(['git', '-C', args.git_repo, 'rev-parse', args.benchmark_commit]).decode().strip() + return info + +def dict_to_params(d): + """ + Return a list of 'key=value' strings from a dictionary. + """ + res = [] + for (k, v) in d.items(): + res.append(f'{k}={v}') + return res + +def main(argv): + parser = argparse.ArgumentParser( + prog='run-benchmarks', + description='Benchmark libc++ at the given commit and submit to LNT.') + parser.add_argument('--benchmark-commit', type=str, required=True, + help='The SHA representing the version of the library to benchmark.') + parser.add_argument('--test-suite-commit', type=str, required=True, + help='The SHA representing the version of the test suite to use for benchmarking.') + parser.add_argument('--compiler', type=str, required=True, + help='Path to the compiler to use for testing.') + parser.add_argument('--machine', type=str, required=True, + help='The name of the machine for reporting LNT results.') + parser.add_argument('--test-suite', type=str, required=True, + help='The name of the test suite for reporting LNT results.') + parser.add_argument('--lnt-url', type=str, required=True, + help='The URL of the LNT instance to submit results to.') + parser.add_argument('--disable-microbenchmarks', action='store_true', + help="Do not run the microbenchmarks, only run SPEC (if possible).") + parser.add_argument('--spec-dir', type=pathlib.Path, required=False, + help='Optional path to a SPEC installation to use for benchmarking.') + parser.add_argument('--git-repo', type=directory_path, default=os.getcwd(), + help='Optional path to the Git repository to use. By default, the current working directory is used.') + parser.add_argument('--dry-run', action='store_true', + help='Only print what would be executed.') + parser.add_argument('-v', '--verbose', action='store_true', + help='Print the output of all subcommands.') + args = parser.parse_args(argv) + + logging.basicConfig(level=logging.INFO) + + do_spec = args.spec_dir is not None + do_micro = not args.disable_microbenchmarks + if not (do_spec or do_micro): + raise ValueError("You must run at least the microbenchmarks or SPEC") + + def run(command, *posargs, enforce_success=True, **kwargs): + command = [str(c) for c in command] + if args.dry_run or args.verbose: + print(f'$ {" ".join(command)}') + + if not args.dry_run: + # If we're running with verbose, print everything, otherwise, hush everything. + if not args.verbose: + if 'stdout' not in kwargs: + kwargs.update({'stdout': subprocess.DEVNULL}) + if 'stderr' not in kwargs: + kwargs.update({'stderr': subprocess.DEVNULL}) + if enforce_success: + subprocess.check_call(command, *posargs, **kwargs) + else: + subprocess.call(command, *posargs, **kwargs) + + with tempfile.TemporaryDirectory() as build_dir: + build_dir = pathlib.Path(build_dir) + + logging.info('Installing LNT') + run(['python', '-m', 'venv', build_dir / '.venv']) + run([build_dir / '.venv/bin/pip', 'install', 'llvm-lnt']) + + logging.info(f'Building libc++ at commit {args.benchmark_commit}') + run([args.git_repo / 'libcxx/utils/build-at-commit', + '--git-repo', args.git_repo, + '--install-dir', build_dir / 'install', + '--commit', args.benchmark_commit, + '--', '-DCMAKE_BUILD_TYPE=RelWithDebInfo', f'-DCMAKE_CXX_COMPILER={args.compiler}']) + + if do_spec: + logging.info(f'Running SPEC benchmarks from {args.test_suite_commit} against libc++ {args.benchmark_commit}') + run([args.git_repo / 'libcxx/utils/test-at-commit', + '--git-repo', args.git_repo, + '--build-dir', build_dir / 'spec', + '--test-suite-commit', args.test_suite_commit, + '--libcxx-installation', build_dir / 'install', + '--', + '-j1', '--time-tests', + '--param', f'compiler={args.compiler}', + '--param', 'optimization=speed', + '--param', 'std=c++17', + '--param', f'spec_dir={args.spec_dir}', + build_dir / 'spec/libcxx/test', + '--filter', 'benchmarks/spec.gen.py'], + enforce_success=False) + with open(build_dir / 'benchmarks.lnt', 'a') as f: + run([args.git_repo / 'libcxx/utils/consolidate-benchmarks', build_dir / 'spec'], stdout=f) + + # TODO: For now, we run only a subset of the benchmarks because running the whole test suite is too slow. + # Run the whole test suite once https://github.com/llvm/llvm-project/issues/173032 is resolved. + if do_micro: + logging.info(f'Running microbenchmarks from {args.test_suite_commit} against libc++ {args.benchmark_commit}') + run([args.git_repo / 'libcxx/utils/test-at-commit', + '--git-repo', args.git_repo, + '--build-dir', build_dir / 'micro', + '--test-suite-commit', args.test_suite_commit, + '--libcxx-installation', build_dir / 'install', + '--', + '-j1', '--time-tests', + '--param', f'compiler={args.compiler}', + '--param', 'optimization=speed', + '--param', 'std=c++26', + build_dir / 'micro/libcxx/test', + '--filter', 'benchmarks/(algorithms|containers|iterators|locale|memory|streams|numeric|utility)'], + enforce_success=False) + with open(build_dir / 'benchmarks.lnt', 'a') as f: + run([args.git_repo / 'libcxx/utils/consolidate-benchmarks', build_dir / 'micro'], stdout=f) + + logging.info('Creating JSON report for LNT') + order = len(subprocess.check_output(['git', '-C', args.git_repo, 'rev-list', args.benchmark_commit]).splitlines()) + importreport = [build_dir / '.venv/bin/lnt', 'importreport', '--order', str(order), '--machine', args.machine] + for arg in dict_to_params(gather_run_information(args)): + importreport += ['--run-info', arg] + for arg in dict_to_params(gather_machine_information(args)): + importreport += ['--machine-info', arg] + importreport += [build_dir / 'benchmarks.lnt', build_dir / 'benchmarks.json'] + run(importreport) + + logging.info(f'Submitting results to {args.lnt_url}') + submission_url = f'{args.lnt_url}/db_default/v4/{args.test_suite}/submitRun' + run([build_dir / '.venv/bin/lnt', 'submit', '--ignore-regressions', submission_url, build_dir / 'benchmarks.json']) + + +if __name__ == '__main__': + main(sys.argv[1:]) diff --git a/libcxx/utils/ci/lnt/schema.yaml b/libcxx/utils/ci/lnt/schema.yaml new file mode 100644 index 0000000000000..7b4e3da42635e --- /dev/null +++ b/libcxx/utils/ci/lnt/schema.yaml @@ -0,0 +1,44 @@ +# Schema definition for the libc++ benchmark suite when running a LNT +# instance. The metrics in this schema correspond to the metrics gathered +# by libc++'s benchmark suite, and the "run fields" correspond to the +# information gathered by the script that runs those benchmarks for +# uploading with LNT. +format_version: '2' +name: libcxx +metrics: +- name: execution_time + type: Real + display_name: Execution Time + unit: seconds + unit_abbrev: s +- name: instructions + type: Real + display_name: Instructions Retired + unit: instructions + unit_abbrev: instr +- name: max_rss + type: Real + display_name: Maximum Resident Set Size + unit: megabytes + unit_abbrev: mb +- name: cycles + type: Real + display_name: Cycles Elapsed + unit: cycles + unit_abbrev: cycles +- name: peak_memory + type: Real + display_name: Peak Memory Footprint + unit: megabytes + unit_abbrev: mb +run_fields: +- name: llvm_project_revision + order: true +- name: commit_info +- name: git_sha +machine_fields: +- name: hardware +- name: os +- name: test_suite_commit +- name: compiler +- name: sdk From 94298f3411ba799620d7fc81ea1466dfa6f450df Mon Sep 17 00:00:00 2001 From: Kseniya Tikhomirova Date: Wed, 11 Feb 2026 20:19:33 +0100 Subject: [PATCH 004/705] [libsycl] Add lit configuration files and basic test (#177407) This PR brings initial version of lit configs for libsycl tests. We plan to add more tests for libsycl integration with the compiler in the future. --------- Signed-off-by: Tikhomirova, Kseniya --- libsycl/CMakeLists.txt | 11 +- libsycl/test/CMakeLists.txt | 26 +++ libsycl/test/README.md | 83 ++++++++ libsycl/test/basic/platform_get_devices.cpp | 125 ++++++++++++ libsycl/test/lit.cfg.py | 211 ++++++++++++++++++++ libsycl/test/lit.site.cfg.py.in | 34 ++++ 6 files changed, 489 insertions(+), 1 deletion(-) create mode 100644 libsycl/test/CMakeLists.txt create mode 100644 libsycl/test/README.md create mode 100644 libsycl/test/basic/platform_get_devices.cpp create mode 100644 libsycl/test/lit.cfg.py create mode 100644 libsycl/test/lit.site.cfg.py.in diff --git a/libsycl/CMakeLists.txt b/libsycl/CMakeLists.txt index f25f51def0cc7..f75216978f69a 100644 --- a/libsycl/CMakeLists.txt +++ b/libsycl/CMakeLists.txt @@ -133,5 +133,14 @@ add_custom_target(libsycl-runtime-libraries ) add_subdirectory(src) - add_subdirectory(tools) + +if(LLVM_INCLUDE_TESTS) + add_subdirectory(test) +endif() + +add_custom_target(libsycl-toolchain ALL + DEPENDS libsycl-runtime-libraries + sycl-ls + COMMENT "Building libsycl toolchain..." +) diff --git a/libsycl/test/CMakeLists.txt b/libsycl/test/CMakeLists.txt new file mode 100644 index 0000000000000..ed5dce5ae2566 --- /dev/null +++ b/libsycl/test/CMakeLists.txt @@ -0,0 +1,26 @@ +cmake_minimum_required(VERSION 3.20.0) + +message("Configuring libsycl tests") + +set(LIBSYCL_CXX_COMPILER "${LLVM_BINARY_DIR}/bin/clang++") +set(LIBSYCL_TEST_COMPILER_OPTIONS "" CACHE STRING + "SYCL compiler options for building libsycl tests") + +configure_lit_site_cfg( + ${CMAKE_CURRENT_SOURCE_DIR}/lit.site.cfg.py.in + ${CMAKE_CURRENT_BINARY_DIR}/lit.site.cfg.py + MAIN_CONFIG + ${CMAKE_CURRENT_SOURCE_DIR}/lit.cfg.py +) + +list(APPEND LIBSYCL_TEST_DEPS + libsycl-toolchain + FileCheck + not +) + +add_lit_testsuite(check-sycl + "Running libsycl tests" + ${CMAKE_CURRENT_BINARY_DIR} + EXCLUDE_FROM_CHECK_ALL + DEPENDS ${LIBSYCL_TEST_DEPS}) diff --git a/libsycl/test/README.md b/libsycl/test/README.md new file mode 100644 index 0000000000000..7ee8899012eeb --- /dev/null +++ b/libsycl/test/README.md @@ -0,0 +1,83 @@ +## Getting Started + +This directory contains `libsycl` tests distributed in subdirectories based on +testing scope. `libsycl` uses LIT to configure and run its tests. + +Please see the [Lit Command Guide](https://llvm.org/docs/CommandGuide/lit.html) +for more information about LIT. + +## Prerequisites + +* Target runtime(s) to execute tests on devices. + TODO: add link to liboffload instruction once they add it. +* Compiler & libsycl. Can be built following these + [instructions](/libsycl/docs/index.rst). + +## Run the tests + +`libsycl` is integrated via `LLVM_ENABLE_RUNTIMES` and is not visible as a top +level target. The same is applicable to tests. To run `check-sycl` tests you need +to prepend `/runtimes/runtimes-bins/` to the paths of all tests. +For example, to run all the `libsycl` tests you can do: +```bash +/bin/llvm-lit /runtimes/runtimes-bins/libsycl/test +``` + +To run an individual test, use the path to it instead. + +If you are using `ninja` as your build system, you can run all the tests in the +`libsycl` testsuite as: + +```bash + ninja -C /runtimes/runtimes-bins check-sycl + ``` + +## CMake parameters + +These parameters can be used to configure tests: + +`LIBSYCL_CXX_COMPILER` - path to compiler to use it for building tests. + +`LIBSYCL_TEST_COMPILER_OPTIONS` - flags to be passed to `LIBSYCL_CXX_COMPILER` when + building libsycl tests. + +`LLVM_LIT` - path to llvm-lit tool. + +## Creating or modifying tests + +### LIT feature checks + +Following features can be passed to LIT via `REQUIRES`, `UNSUPPORTED`, etc. +filters to limit test execution to the specific environment. + +#### Auto-detected features + +The following features are automatically detected by `llvm-lit` by scanning the +environment: + +* `any-device` - presence of at least 1 supported device; +* `any-device-is-gpu` - device type to be available; +* `any-device-is-level_zero` - backend to be available; + +Note: the `sycl-ls` tool doesn't have an assigned feature since it is essential for +test configuration and is always available. + +### llvm-lit parameters + +The following options can be passed to the `llvm-lit` tool with `--param` +option to configure test execution: + +* `libsycl_compiler` - full path to compiler to use. +* `extra_environment` - comma-separated list of variables with values to be + added to the test environment. Can be also set by the `LIT_EXTRA_ENVIRONMENT` + variable in CMake. +* `extra_system_environment` - comma-separated list of variables to be + propagated from the host environment to the test environment. Can be also set + by `LIT_EXTRA_SYSTEM_ENVIRONMENT` variable in CMake. + +Example: + +```bash +/bin/llvm-lit --param libsycl_compiler=path/to/clang++ \ + /runtimes/runtimes-bins/libsycl/test +``` \ No newline at end of file diff --git a/libsycl/test/basic/platform_get_devices.cpp b/libsycl/test/basic/platform_get_devices.cpp new file mode 100644 index 0000000000000..ea8652a313e8a --- /dev/null +++ b/libsycl/test/basic/platform_get_devices.cpp @@ -0,0 +1,125 @@ +// REQUIRES: any-device +// RUN: %clangxx %sycl_options %s -o %t.out +// RUN: %t.out +// +// Tests platform::get_devices for each device type. + +#include + +#include +#include + +std::string BackendToString(sycl::backend Backend) { + switch (Backend) { + case sycl::backend::opencl: + return "opencl"; + case sycl::backend::level_zero: + return "level_zero"; + case sycl::backend::cuda: + return "cuda"; + case sycl::backend::hip: + return "hip"; + default: + return "unknown"; + } +} + +std::string DeviceTypeToString(sycl::info::device_type DevType) { + switch (DevType) { + case sycl::info::device_type::all: + return "device_type::all"; + case sycl::info::device_type::cpu: + return "device_type::cpu"; + case sycl::info::device_type::gpu: + return "device_type::gpu"; + case sycl::info::device_type::accelerator: + return "device_type::accelerator"; + case sycl::info::device_type::custom: + return "device_type::custom"; + case sycl::info::device_type::automatic: + return "device_type::automatic"; + case sycl::info::device_type::host: + return "device_type::host"; + default: + return "unknown"; + } +} + +std::string GenerateDeviceDescription(sycl::info::device_type DevType, + const sycl::platform &Platform) { + return std::string(DeviceTypeToString(DevType)) + " (" + + BackendToString(Platform.get_backend()) + ")"; +} + +template +int Check(const T1 &LHS, const T2 &RHS, std::string TestName) { + if (LHS == RHS) + return 0; + + std::cerr << "Failed check " << LHS << " != " << RHS << ": " << TestName + << std::endl; + return 1; +} + +int CheckDeviceType(const sycl::platform &P, sycl::info::device_type DevType, + std::vector &AllDevices) { + // This check verifies data of device with specific device_type and if it is + // correctly chosen among all devices (device_type::all). + // Though there is no point to check device_type::all here. + assert(DevType != sycl::info::device_type::all); + int Failures = 0; + + std::vector Devices = P.get_devices(DevType); + + if (DevType == sycl::info::device_type::automatic) { + if (AllDevices.empty()) { + Failures += Check(Devices.size(), 0, + "No devices reported for device_type::all query, but " + "device_type::automatic returns a device."); + } else { + Failures += Check(Devices.size(), 1, + "Number of devices for device_type::automatic query."); + if (Devices.size()) + Failures += Check( + std::count(AllDevices.begin(), AllDevices.end(), Devices[0]), 1, + "Device is in the set of device_type::all devices in the " + "platform."); + } + return Failures; + } + + // Count devices with the type. + size_t DevCount = 0; + for (sycl::device Device : Devices) + DevCount += (Device.get_info() == DevType); + + Failures += Check(Devices.size(), DevCount, + "Unexpected number of devices for " + + GenerateDeviceDescription(DevType, P)); + + Failures += + Check(std::all_of(Devices.begin(), Devices.end(), + [&](const auto &Dev) { + return std::count(AllDevices.begin(), + AllDevices.end(), Dev) == 1; + }), + true, + "Not all devices for " + GenerateDeviceDescription(DevType, P) + + " appear in the list of all devices"); + + return Failures; +} + +int main() { + int Failures = 0; + for (sycl::platform P : sycl::platform::get_platforms()) { + std::vector Devices = P.get_devices(); + + for (sycl::info::device_type DevType : + {sycl::info::device_type::cpu, sycl::info::device_type::gpu, + sycl::info::device_type::accelerator, sycl::info::device_type::custom, + sycl::info::device_type::automatic, sycl::info::device_type::host}) + Failures += CheckDeviceType(P, DevType, Devices); + } + return Failures; +} diff --git a/libsycl/test/lit.cfg.py b/libsycl/test/lit.cfg.py new file mode 100644 index 0000000000000..67ab5d2e83294 --- /dev/null +++ b/libsycl/test/lit.cfg.py @@ -0,0 +1,211 @@ +# -*- Python -*- + +# Configuration file for the 'lit' test runner. + +import os +import subprocess + +from lit.llvm import llvm_config +import lit.formats +from lit.llvm.subst import ToolSubst, FindTool + +# name: The name of this test suite. +config.name = "libsycl" + +# suffixes: A list of file extensions to treat as test files. +config.suffixes = [".cpp"] + +config.excludes = ["Inputs"] + +# test_source_root: The root path where tests are located. +config.test_source_root = os.path.dirname(__file__) + +# allow expanding substitutions that are based on other substitutions +config.recursiveExpansionLimit = 10 + +# test_exec_root: The root path where tests should be run. +config.test_exec_root = config.libsycl_obj_root + +# To be filled by lit.local.cfg files. +config.required_features = [] +config.unsupported_features = [] + +# Cleanup environment variables which may affect tests +possibly_dangerous_env_vars = [ + "COMPILER_PATH", + "RC_DEBUG_OPTIONS", + "CINDEXTEST_PREAMBLE_FILE", + "LIBRARY_PATH", + "CPATH", + "C_INCLUDE_PATH", + "CPLUS_INCLUDE_PATH", + "OBJC_INCLUDE_PATH", + "OBJCPLUS_INCLUDE_PATH", + "LIBCLANG_TIMING", + "LIBCLANG_OBJTRACKING", + "LIBCLANG_LOGGING", + "LIBCLANG_BGPRIO_INDEX", + "LIBCLANG_BGPRIO_EDIT", + "LIBCLANG_NOTHREADS", + "LIBCLANG_RESOURCE_USAGE", + "LIBCLANG_CODE_COMPLETION_LOGGING", + "INCLUDE", +] + +for name in possibly_dangerous_env_vars: + if name in llvm_config.config.environment: + del llvm_config.config.environment[name] + +# Propagate some variables from the host environment. +llvm_config.with_system_environment( + [ + "PATH", + ] +) + +# Take into account extra system environment variables if provided via parameter. +if config.extra_system_environment: + lit_config.note( + "Extra system variables to propagate value from: " + + config.extra_system_environment + ) + extra_env_vars = config.extra_system_environment.split(",") + for var in extra_env_vars: + if var in os.environ: + llvm_config.with_system_environment(var) + +llvm_config.with_environment("PATH", config.lit_tools_dir, append_path=True) + +# Configure LD_LIBRARY_PATH +llvm_config.with_system_environment( + ["LD_LIBRARY_PATH", "LIBRARY_PATH", "C_INCLUDE_PATH", "CPLUS_INCLUDE_PATH"] +) +llvm_config.with_environment( + "LD_LIBRARY_PATH", config.libsycl_libs_dir, append_path=True +) + +llvm_config.with_environment("PATH", config.libsycl_tools_dir, append_path=True) + +if config.extra_environment: + lit_config.note("Extra environment variables") + for env_pair in config.extra_environment.split(","): + [var, val] = env_pair.split("=", 1) + if val: + llvm_config.with_environment(var, val) + lit_config.note("\t" + var + "=" + val) + else: + lit_config.note("\tUnset " + var) + llvm_config.with_environment(var, "") + + +# Temporarily modify environment to be the same that we use when running tests +class test_env: + def __enter__(self): + self.old_environ = dict(os.environ) + os.environ.clear() + os.environ.update(config.environment) + self.old_dir = os.getcwd() + os.chdir(config.libsycl_obj_root) + + def __exit__(self, exc_type, exc_value, exc_traceback): + os.environ.clear() + os.environ.update(self.old_environ) + os.chdir(self.old_dir) + + +# General substritutions +config.substitutions.append( + ( + "%sycl_options", + " -lsycl" + + " -isystem " + + config.libsycl_include + + " -isystem " + + os.path.join(config.libsycl_include, "sycl") + + " -L" + + config.libsycl_libs_dir, + ) +) +config.substitutions.append(("%sycl_libs_dir", config.libsycl_libs_dir)) +config.substitutions.append(("%sycl_static_libs_dir", config.libsycl_libs_dir)) +config.substitutions.append(("%obj_ext", ".o")) +config.substitutions.append(("%sycl_include", "-isystem " + config.libsycl_include)) +config.substitutions.append(("%include_option", "-include")) +config.substitutions.append(("%debug_option", "-g")) +config.substitutions.append(("%cxx_std_option", "-std=")) +config.substitutions.append(("%fPIC", "-fPIC")) +config.substitutions.append(("%shared_lib", "-shared")) +config.substitutions.append(("%O0", "-O0")) + +sycl_ls = FindTool("sycl-ls").resolve( + llvm_config, os.pathsep.join([config.libsycl_bin_dir, config.llvm_tools_dir]) +) +if not sycl_ls: + lit_config.fatal("can't find `sycl-ls`") + +tools = [ + ToolSubst("FileCheck", unresolved="ignore"), + # not is only substituted in certain circumstances; this is lit's default + # behaviour. + ToolSubst( + r"\| \bnot\b", command=FindTool("not"), verbatim=True, unresolved="ignore" + ), + ToolSubst("sycl-ls", command=sycl_ls, unresolved="fatal"), +] + +# Try and find each of these tools in the libsycl bin directory, in the llvm tools directory +# or the PATH, in that order. If found, they will be added as substitutions with the full path +# to the tool. +llvm_config.add_tool_substitutions( + tools, [config.libsycl_bin_dir, config.llvm_tools_dir, os.environ.get("PATH", "")] +) + +lit_config.note("Targeted devices: all") +with test_env(): + sycl_ls_output = subprocess.check_output(sycl_ls, text=True, shell=True) + + devices = set() + for line in sycl_ls_output.splitlines(): + if not line.startswith("["): + continue + backend, device = line[1:].split("]")[0].split(":") + devices.add("{}:{}".format(backend, device)) + libsycl_devices = list(devices) + +available_devices = { + "level_zero": "gpu", +} +for d in libsycl_devices: + be, dev = d.split(":") + if be not in available_devices: + lit_config.error("Unsupported device {}".format(d)) + if dev not in available_devices[be]: + lit_config.error("Unsupported device {}".format(d)) + +if len(libsycl_devices) > 0: + config.available_features.add("any-device") + +for sycl_device in libsycl_devices: + be, dev = sycl_device.split(":") + config.available_features.add("any-device-is-" + dev) + config.available_features.add("any-device-is-" + be) + +# Check if user passed verbose-print parameter, if yes, add VERBOSE_PRINT macro +if "verbose-print" in lit_config.params: + verbose_print = "-DVERBOSE_PRINT" +else: + verbose_print = "" + +clangxx = " " + config.libsycl_compiler + " -Werror " + config.cxx_flags + verbose_print +config.substitutions.append(("%clangxx", clangxx)) + +config.test_format = lit.formats.ShTest() + +try: + import psutil + + # Set timeout for a single test + lit_config.maxIndividualTestTime = 600 + +except ImportError: + pass diff --git a/libsycl/test/lit.site.cfg.py.in b/libsycl/test/lit.site.cfg.py.in new file mode 100644 index 0000000000000..a0d9c60b3e4b9 --- /dev/null +++ b/libsycl/test/lit.site.cfg.py.in @@ -0,0 +1,34 @@ +@LIT_SITE_CFG_IN_HEADER@ + +import subprocess +import site + +site.addsitedir("@CMAKE_CURRENT_SOURCE_DIR@") + +config.libsycl_compiler = lit_config.params.get("libsycl_compiler", "@LIBSYCL_CXX_COMPILER@") +config.libsycl_bin_dir = os.path.dirname(config.libsycl_compiler) +config.libsycl_root_dir = os.path.dirname(config.libsycl_bin_dir) + +config.cxx_flags = lit_config.params.get("cxx_flags", "@LIBSYCL_TEST_COMPILER_OPTIONS@") + +config.extra_environment = lit_config.params.get("extra_environment", "@LIT_EXTRA_ENVIRONMENT@") +config.extra_system_environment = lit_config.params.get("extra_system_environment", "@LIT_EXTRA_SYSTEM_ENVIRONMENT@") + +def get_libsycl_tool_path(name): + try: + return subprocess.check_output([config.libsycl_compiler, "-print-prog-name=" + name], text=True) + except subprocess.CalledProcessError: + return os.path.join(config.libsycl_bin_dir, name) + +config.llvm_tools_dir = os.path.dirname(get_libsycl_tool_path("llvm-config")) +config.lit_tools_dir = os.path.dirname("@TEST_SUITE_LIT@") + +config.libsycl_tools_dir = config.llvm_tools_dir +config.libsycl_include = os.path.join(config.libsycl_root_dir, 'include') +config.libsycl_obj_root = "@CMAKE_CURRENT_BINARY_DIR@" +config.libsycl_libs_dir = os.path.join(config.libsycl_root_dir, 'lib/x86_64-unknown-linux-gnu') + +import lit.llvm +lit.llvm.initialize(lit_config, config) + +lit_config.load_config(config, "@CMAKE_CURRENT_SOURCE_DIR@/lit.cfg.py") From ee722ae1df203b0e246ea2b68315c6b303edfcf6 Mon Sep 17 00:00:00 2001 From: Florian Hahn Date: Wed, 11 Feb 2026 19:21:28 +0000 Subject: [PATCH 005/705] [VPlan] Remove VPUnrollPartAccessor from VPReductionPHIRecipe (NFC). VPUnrollPartAccessor is not used any longer by VPReductionPHIRecipe. Remove it. --- llvm/lib/Transforms/Vectorize/VPlan.h | 3 +-- llvm/lib/Transforms/Vectorize/VPlanUnroll.cpp | 1 - 2 files changed, 1 insertion(+), 3 deletions(-) diff --git a/llvm/lib/Transforms/Vectorize/VPlan.h b/llvm/lib/Transforms/Vectorize/VPlan.h index eac5b58841e80..68ff4e6a037f5 100644 --- a/llvm/lib/Transforms/Vectorize/VPlan.h +++ b/llvm/lib/Transforms/Vectorize/VPlan.h @@ -2598,8 +2598,7 @@ inline ReductionStyle getReductionStyle(bool InLoop, bool Ordered, /// A recipe for handling reduction phis. The start value is the first operand /// of the recipe and the incoming value from the backedge is the second /// operand. -class VPReductionPHIRecipe : public VPHeaderPHIRecipe, - public VPUnrollPartAccessor<2> { +class VPReductionPHIRecipe : public VPHeaderPHIRecipe { /// The recurrence kind of the reduction. const RecurKind Kind; diff --git a/llvm/lib/Transforms/Vectorize/VPlanUnroll.cpp b/llvm/lib/Transforms/Vectorize/VPlanUnroll.cpp index 53cac9fcd80d6..d2b5583f06c3f 100644 --- a/llvm/lib/Transforms/Vectorize/VPlanUnroll.cpp +++ b/llvm/lib/Transforms/Vectorize/VPlanUnroll.cpp @@ -284,7 +284,6 @@ void UnrollState::unrollHeaderPHIByUF(VPHeaderPHIRecipe *R, for (unsigned Part = 1; Part != UF; ++Part) VPV2Parts[VPI][Part - 1] = StartV; } - Copy->addOperand(getConstantInt(Part)); } else { assert(isa(R) && "unexpected header phi recipe not needing unrolled part"); From 65f3329a556c5b6f0261ff9268c776a5c7864088 Mon Sep 17 00:00:00 2001 From: Syadus Sefat <42645939+mssefat@users.noreply.github.com> Date: Wed, 11 Feb 2026 13:24:08 -0600 Subject: [PATCH 006/705] [AMDGPU][GlobalIsel] Add register bank legalization rules for buffer atomic int min max (#180975) This patch adds register bank legalization rules for buffer atomic int min max operations in the AMDGPU GlobalISel pipeline. --- .../AMDGPU/AMDGPURegBankLegalizeRules.cpp | 5 +- ...amdgcn.raw.buffer.atomic.integer-minmax.ll | 282 ++++++++++++++++++ 2 files changed, 286 insertions(+), 1 deletion(-) create mode 100644 llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.raw.buffer.atomic.integer-minmax.ll diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp index 6f4be08beee4e..a0be07d7eae05 100644 --- a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp +++ b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp @@ -817,7 +817,10 @@ RegBankLegalizeRules::RegBankLegalizeRules(const GCNSubtarget &_ST, .Div(S64, {{Vgpr64}, {Vgpr64, Vgpr64, SgprV4S32_WF, Vgpr32, Vgpr32, Sgpr32_WF}}); - addRulesForGOpcs({G_AMDGPU_BUFFER_ATOMIC_SWAP}, Standard) + addRulesForGOpcs({G_AMDGPU_BUFFER_ATOMIC_SWAP, G_AMDGPU_BUFFER_ATOMIC_UMAX, + G_AMDGPU_BUFFER_ATOMIC_UMIN, G_AMDGPU_BUFFER_ATOMIC_SMAX, + G_AMDGPU_BUFFER_ATOMIC_SMIN}, + Standard) .Div(S32, {{Vgpr32}, {Vgpr32, SgprV4S32_WF, Vgpr32, Vgpr32, Sgpr32_WF}}) .Div(S64, {{Vgpr64}, {Vgpr64, SgprV4S32_WF, Vgpr32, Vgpr32, Sgpr32_WF}}); diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.raw.buffer.atomic.integer-minmax.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.raw.buffer.atomic.integer-minmax.ll new file mode 100644 index 0000000000000..d3cd8839d127b --- /dev/null +++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/llvm.amdgcn.raw.buffer.atomic.integer-minmax.ll @@ -0,0 +1,282 @@ +; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6 +; RUN: llc -debug -global-isel -mtriple=amdgcn -mcpu=gfx1010 -new-reg-bank-select < %s | FileCheck %s + +define amdgpu_ps float @test1(i32 %val, <4 x i32> inreg %rsrc, i32 %voffset, i32 inreg %soffset) { +; CHECK-LABEL: test1: +; CHECK: ; %bb.0: +; CHECK-NEXT: buffer_atomic_smin v0, v1, s[0:3], s4 offen glc +; CHECK-NEXT: s_waitcnt vmcnt(0) +; CHECK-NEXT: ; return to shader part epilog + %ret = call i32 @llvm.amdgcn.raw.buffer.atomic.smin.i32(i32 %val, <4 x i32> %rsrc, i32 %voffset, i32 %soffset, i32 0) + %cast = bitcast i32 %ret to float + ret float %cast +} + +define amdgpu_ps float @test2(i32 %val, <4 x i32> inreg %rsrc, i32 %voffset, i32 inreg %soffset) { +; CHECK-LABEL: test2: +; CHECK: ; %bb.0: +; CHECK-NEXT: buffer_atomic_smax v0, v1, s[0:3], s4 offen glc +; CHECK-NEXT: s_waitcnt vmcnt(0) +; CHECK-NEXT: ; return to shader part epilog + %ret = call i32 @llvm.amdgcn.raw.buffer.atomic.smax.i32(i32 %val, <4 x i32> %rsrc, i32 %voffset, i32 %soffset, i32 0) + %cast = bitcast i32 %ret to float + ret float %cast +} + +define amdgpu_ps float @test3(i32 %val, <4 x i32> inreg %rsrc, i32 %voffset, i32 inreg %soffset) { +; CHECK-LABEL: test3: +; CHECK: ; %bb.0: +; CHECK-NEXT: buffer_atomic_umin v0, v1, s[0:3], s4 offen glc +; CHECK-NEXT: s_waitcnt vmcnt(0) +; CHECK-NEXT: ; return to shader part epilog + %ret = call i32 @llvm.amdgcn.raw.buffer.atomic.umin.i32(i32 %val, <4 x i32> %rsrc, i32 %voffset, i32 %soffset, i32 0) + %cast = bitcast i32 %ret to float + ret float %cast +} + +define amdgpu_ps float @test4(i32 %val, <4 x i32> inreg %rsrc, i32 %voffset, i32 inreg %soffset) { +; CHECK-LABEL: test4: +; CHECK: ; %bb.0: +; CHECK-NEXT: buffer_atomic_umax v0, v1, s[0:3], s4 offen glc +; CHECK-NEXT: s_waitcnt vmcnt(0) +; CHECK-NEXT: ; return to shader part epilog + %ret = call i32 @llvm.amdgcn.raw.buffer.atomic.umax.i32(i32 %val, <4 x i32> %rsrc, i32 %voffset, i32 %soffset, i32 0) + %cast = bitcast i32 %ret to float + ret float %cast +} + + +define amdgpu_ps <2 x float> @test5(i64 %val, <4 x i32> inreg %rsrc, i32 %voffset, i32 inreg %soffset) { +; CHECK-LABEL: test5: +; CHECK: ; %bb.0: +; CHECK-NEXT: buffer_atomic_smin_x2 v[0:1], v2, s[0:3], s4 offen glc +; CHECK-NEXT: s_waitcnt vmcnt(0) +; CHECK-NEXT: ; return to shader part epilog + %ret = call i64 @llvm.amdgcn.raw.buffer.atomic.smin.i64(i64 %val, <4 x i32> %rsrc, i32 %voffset, i32 %soffset, i32 0) + %cast = bitcast i64 %ret to <2 x float> + ret <2 x float> %cast +} + +define amdgpu_ps <2 x float> @test6(i64 %val, <4 x i32> inreg %rsrc, i32 %voffset, i32 inreg %soffset) { +; CHECK-LABEL: test6: +; CHECK: ; %bb.0: +; CHECK-NEXT: buffer_atomic_smax_x2 v[0:1], v2, s[0:3], s4 offen glc +; CHECK-NEXT: s_waitcnt vmcnt(0) +; CHECK-NEXT: ; return to shader part epilog + %ret = call i64 @llvm.amdgcn.raw.buffer.atomic.smax.i64(i64 %val, <4 x i32> %rsrc, i32 %voffset, i32 %soffset, i32 0) + %cast = bitcast i64 %ret to <2 x float> + ret <2 x float> %cast +} + +define amdgpu_ps <2 x float> @test7(i64 %val, <4 x i32> inreg %rsrc, i32 %voffset, i32 inreg %soffset) { +; CHECK-LABEL: test7: +; CHECK: ; %bb.0: +; CHECK-NEXT: buffer_atomic_umin_x2 v[0:1], v2, s[0:3], s4 offen glc +; CHECK-NEXT: s_waitcnt vmcnt(0) +; CHECK-NEXT: ; return to shader part epilog + %ret = call i64 @llvm.amdgcn.raw.buffer.atomic.umin.i64(i64 %val, <4 x i32> %rsrc, i32 %voffset, i32 %soffset, i32 0) + %cast = bitcast i64 %ret to <2 x float> + ret <2 x float> %cast +} + +define amdgpu_ps <2 x float> @test8(i64 %val, <4 x i32> inreg %rsrc, i32 %voffset, i32 inreg %soffset) { +; CHECK-LABEL: test8: +; CHECK: ; %bb.0: +; CHECK-NEXT: buffer_atomic_umax_x2 v[0:1], v2, s[0:3], s4 offen glc +; CHECK-NEXT: s_waitcnt vmcnt(0) +; CHECK-NEXT: ; return to shader part epilog + %ret = call i64 @llvm.amdgcn.raw.buffer.atomic.umax.i64(i64 %val, <4 x i32> %rsrc, i32 %voffset, i32 %soffset, i32 0) + %cast = bitcast i64 %ret to <2 x float> + ret <2 x float> %cast +} + +define amdgpu_ps float @wf_test1(i32 %val, <4 x i32> %rsrc, i32 %voffset, i32 inreg %soffset) { +; CHECK-LABEL: wf_test1: +; CHECK: ; %bb.0: +; CHECK-NEXT: s_mov_b32 s2, exec_lo +; CHECK-NEXT: .LBB8_1: ; =>This Inner Loop Header: Depth=1 +; CHECK-NEXT: v_readfirstlane_b32 s4, v1 +; CHECK-NEXT: v_readfirstlane_b32 s5, v2 +; CHECK-NEXT: v_readfirstlane_b32 s6, v3 +; CHECK-NEXT: v_readfirstlane_b32 s7, v4 +; CHECK-NEXT: v_cmp_eq_u64_e32 vcc_lo, s[4:5], v[1:2] +; CHECK-NEXT: v_cmp_eq_u64_e64 s1, s[6:7], v[3:4] +; CHECK-NEXT: s_and_b32 s1, vcc_lo, s1 +; CHECK-NEXT: s_and_saveexec_b32 s1, s1 +; CHECK-NEXT: s_waitcnt vmcnt(0) +; CHECK-NEXT: buffer_atomic_smin v0, v5, s[4:7], s0 offen glc +; CHECK-NEXT: ; implicit-def: $vgpr1 +; CHECK-NEXT: ; implicit-def: $vgpr5 +; CHECK-NEXT: ; implicit-def: $vgpr3_vgpr4 +; CHECK-NEXT: s_waitcnt_depctr depctr_vm_vsrc(0) +; CHECK-NEXT: s_xor_b32 exec_lo, exec_lo, s1 +; CHECK-NEXT: s_cbranch_execnz .LBB8_1 +; CHECK-NEXT: ; %bb.2: +; CHECK-NEXT: s_mov_b32 exec_lo, s2 +; CHECK-NEXT: s_waitcnt vmcnt(0) +; CHECK-NEXT: ; return to shader part epilog + %ret = call i32 @llvm.amdgcn.raw.buffer.atomic.smin.i32(i32 %val, <4 x i32> %rsrc, i32 %voffset, i32 %soffset, i32 0) + %cast = bitcast i32 %ret to float + ret float %cast +} + +define amdgpu_ps <2 x float> @wf_test2(i64 %val, <4 x i32> %rsrc, i32 %voffset, i32 inreg %soffset) { +; CHECK-LABEL: wf_test2: +; CHECK: ; %bb.0: +; CHECK-NEXT: s_mov_b32 s2, exec_lo +; CHECK-NEXT: .LBB9_1: ; =>This Inner Loop Header: Depth=1 +; CHECK-NEXT: v_readfirstlane_b32 s4, v2 +; CHECK-NEXT: v_readfirstlane_b32 s5, v3 +; CHECK-NEXT: v_readfirstlane_b32 s6, v4 +; CHECK-NEXT: v_readfirstlane_b32 s7, v5 +; CHECK-NEXT: v_cmp_eq_u64_e32 vcc_lo, s[4:5], v[2:3] +; CHECK-NEXT: v_cmp_eq_u64_e64 s1, s[6:7], v[4:5] +; CHECK-NEXT: s_and_b32 s1, vcc_lo, s1 +; CHECK-NEXT: s_and_saveexec_b32 s1, s1 +; CHECK-NEXT: s_waitcnt vmcnt(0) +; CHECK-NEXT: buffer_atomic_smax_x2 v[0:1], v6, s[4:7], s0 offen glc +; CHECK-NEXT: ; implicit-def: $vgpr2 +; CHECK-NEXT: ; implicit-def: $vgpr6 +; CHECK-NEXT: ; implicit-def: $vgpr4_vgpr5 +; CHECK-NEXT: s_waitcnt_depctr depctr_vm_vsrc(0) +; CHECK-NEXT: s_xor_b32 exec_lo, exec_lo, s1 +; CHECK-NEXT: s_cbranch_execnz .LBB9_1 +; CHECK-NEXT: ; %bb.2: +; CHECK-NEXT: s_mov_b32 exec_lo, s2 +; CHECK-NEXT: s_waitcnt vmcnt(0) +; CHECK-NEXT: ; return to shader part epilog + %ret = call i64 @llvm.amdgcn.raw.buffer.atomic.smax.i64(i64 %val, <4 x i32> %rsrc, i32 %voffset, i32 %soffset, i32 0) + %cast = bitcast i64 %ret to <2 x float> + ret <2 x float> %cast +} + +; Waterfall tests - divergent soffset (VGPR soffset) +define amdgpu_ps float @wf_test3(i32 %val, <4 x i32> inreg %rsrc, i32 %voffset, i32 %soffset) { +; CHECK-LABEL: wf_test3: +; CHECK: ; %bb.0: +; CHECK-NEXT: s_mov_b32 s4, exec_lo +; CHECK-NEXT: .LBB10_1: ; =>This Inner Loop Header: Depth=1 +; CHECK-NEXT: v_readfirstlane_b32 s6, v2 +; CHECK-NEXT: v_cmp_eq_u32_e32 vcc_lo, s6, v2 +; CHECK-NEXT: s_and_saveexec_b32 s5, vcc_lo +; CHECK-NEXT: s_waitcnt vmcnt(0) +; CHECK-NEXT: buffer_atomic_umin v0, v1, s[0:3], s6 offen glc +; CHECK-NEXT: ; implicit-def: $vgpr2 +; CHECK-NEXT: ; implicit-def: $vgpr1 +; CHECK-NEXT: s_waitcnt_depctr depctr_vm_vsrc(0) +; CHECK-NEXT: s_xor_b32 exec_lo, exec_lo, s5 +; CHECK-NEXT: s_cbranch_execnz .LBB10_1 +; CHECK-NEXT: ; %bb.2: +; CHECK-NEXT: s_mov_b32 exec_lo, s4 +; CHECK-NEXT: s_waitcnt vmcnt(0) +; CHECK-NEXT: ; return to shader part epilog + %ret = call i32 @llvm.amdgcn.raw.buffer.atomic.umin.i32(i32 %val, <4 x i32> %rsrc, i32 %voffset, i32 %soffset, i32 0) + %cast = bitcast i32 %ret to float + ret float %cast +} + +define amdgpu_ps <2 x float> @wf_test5(i64 %val, <4 x i32> inreg %rsrc, i32 %voffset, i32 %soffset) { +; CHECK-LABEL: wf_test5: +; CHECK: ; %bb.0: +; CHECK-NEXT: s_mov_b32 s4, exec_lo +; CHECK-NEXT: .LBB11_1: ; =>This Inner Loop Header: Depth=1 +; CHECK-NEXT: v_readfirstlane_b32 s6, v3 +; CHECK-NEXT: v_cmp_eq_u32_e32 vcc_lo, s6, v3 +; CHECK-NEXT: s_and_saveexec_b32 s5, vcc_lo +; CHECK-NEXT: s_waitcnt vmcnt(0) +; CHECK-NEXT: buffer_atomic_umax_x2 v[0:1], v2, s[0:3], s6 offen glc +; CHECK-NEXT: ; implicit-def: $vgpr3 +; CHECK-NEXT: ; implicit-def: $vgpr2 +; CHECK-NEXT: s_waitcnt_depctr depctr_vm_vsrc(0) +; CHECK-NEXT: s_xor_b32 exec_lo, exec_lo, s5 +; CHECK-NEXT: s_cbranch_execnz .LBB11_1 +; CHECK-NEXT: ; %bb.2: +; CHECK-NEXT: s_mov_b32 exec_lo, s4 +; CHECK-NEXT: s_waitcnt vmcnt(0) +; CHECK-NEXT: ; return to shader part epilog + %ret = call i64 @llvm.amdgcn.raw.buffer.atomic.umax.i64(i64 %val, <4 x i32> %rsrc, i32 %voffset, i32 %soffset, i32 0) + %cast = bitcast i64 %ret to <2 x float> + ret <2 x float> %cast +} + +; Waterfall tests - both divergent (VGPR rsrc and soffset) +define amdgpu_ps float @wf_test6(i32 %val, <4 x i32> %rsrc, i32 %voffset, i32 %soffset) { +; CHECK-LABEL: wf_test6: +; CHECK: ; %bb.0: +; CHECK-NEXT: s_mov_b32 s2, exec_lo +; CHECK-NEXT: .LBB12_1: ; =>This Inner Loop Header: Depth=1 +; CHECK-NEXT: v_readfirstlane_b32 s4, v1 +; CHECK-NEXT: v_readfirstlane_b32 s5, v2 +; CHECK-NEXT: v_readfirstlane_b32 s6, v3 +; CHECK-NEXT: v_readfirstlane_b32 s7, v4 +; CHECK-NEXT: v_readfirstlane_b32 s3, v6 +; CHECK-NEXT: v_cmp_eq_u64_e32 vcc_lo, s[4:5], v[1:2] +; CHECK-NEXT: v_cmp_eq_u64_e64 s0, s[6:7], v[3:4] +; CHECK-NEXT: v_cmp_eq_u32_e64 s1, s3, v6 +; CHECK-NEXT: s_and_b32 s0, vcc_lo, s0 +; CHECK-NEXT: s_and_b32 s0, s0, s1 +; CHECK-NEXT: s_and_saveexec_b32 s0, s0 +; CHECK-NEXT: s_waitcnt vmcnt(0) +; CHECK-NEXT: buffer_atomic_smin v0, v5, s[4:7], s3 offen glc +; CHECK-NEXT: ; implicit-def: $vgpr1 +; CHECK-NEXT: ; implicit-def: $vgpr6 +; CHECK-NEXT: ; implicit-def: $vgpr5 +; CHECK-NEXT: ; implicit-def: $vgpr3_vgpr4 +; CHECK-NEXT: s_waitcnt_depctr depctr_vm_vsrc(0) +; CHECK-NEXT: s_xor_b32 exec_lo, exec_lo, s0 +; CHECK-NEXT: s_cbranch_execnz .LBB12_1 +; CHECK-NEXT: ; %bb.2: +; CHECK-NEXT: s_mov_b32 exec_lo, s2 +; CHECK-NEXT: s_waitcnt vmcnt(0) +; CHECK-NEXT: ; return to shader part epilog + %ret = call i32 @llvm.amdgcn.raw.buffer.atomic.smin.i32(i32 %val, <4 x i32> %rsrc, i32 %voffset, i32 %soffset, i32 0) + %cast = bitcast i32 %ret to float + ret float %cast +} + +define amdgpu_ps <2 x float> @wf_test7(i64 %val, <4 x i32> %rsrc, i32 %voffset, i32 %soffset) { +; CHECK-LABEL: wf_test7: +; CHECK: ; %bb.0: +; CHECK-NEXT: s_mov_b32 s2, exec_lo +; CHECK-NEXT: .LBB13_1: ; =>This Inner Loop Header: Depth=1 +; CHECK-NEXT: v_readfirstlane_b32 s4, v2 +; CHECK-NEXT: v_readfirstlane_b32 s5, v3 +; CHECK-NEXT: v_readfirstlane_b32 s6, v4 +; CHECK-NEXT: v_readfirstlane_b32 s7, v5 +; CHECK-NEXT: v_readfirstlane_b32 s3, v7 +; CHECK-NEXT: v_cmp_eq_u64_e32 vcc_lo, s[4:5], v[2:3] +; CHECK-NEXT: v_cmp_eq_u64_e64 s0, s[6:7], v[4:5] +; CHECK-NEXT: v_cmp_eq_u32_e64 s1, s3, v7 +; CHECK-NEXT: s_and_b32 s0, vcc_lo, s0 +; CHECK-NEXT: s_and_b32 s0, s0, s1 +; CHECK-NEXT: s_and_saveexec_b32 s0, s0 +; CHECK-NEXT: s_waitcnt vmcnt(0) +; CHECK-NEXT: buffer_atomic_umin_x2 v[0:1], v6, s[4:7], s3 offen glc +; CHECK-NEXT: ; implicit-def: $vgpr2 +; CHECK-NEXT: ; implicit-def: $vgpr7 +; CHECK-NEXT: ; implicit-def: $vgpr6 +; CHECK-NEXT: ; implicit-def: $vgpr4_vgpr5 +; CHECK-NEXT: s_waitcnt_depctr depctr_vm_vsrc(0) +; CHECK-NEXT: s_xor_b32 exec_lo, exec_lo, s0 +; CHECK-NEXT: s_cbranch_execnz .LBB13_1 +; CHECK-NEXT: ; %bb.2: +; CHECK-NEXT: s_mov_b32 exec_lo, s2 +; CHECK-NEXT: s_waitcnt vmcnt(0) +; CHECK-NEXT: ; return to shader part epilog + %ret = call i64 @llvm.amdgcn.raw.buffer.atomic.umin.i64(i64 %val, <4 x i32> %rsrc, i32 %voffset, i32 %soffset, i32 0) + %cast = bitcast i64 %ret to <2 x float> + ret <2 x float> %cast +} + + +declare i32 @llvm.amdgcn.raw.buffer.atomic.smin.i32(i32, <4 x i32>, i32, i32, i32 immarg) #0 +declare i32 @llvm.amdgcn.raw.buffer.atomic.smax.i32(i32, <4 x i32>, i32, i32, i32 immarg) #0 +declare i32 @llvm.amdgcn.raw.buffer.atomic.umin.i32(i32, <4 x i32>, i32, i32, i32 immarg) #0 +declare i32 @llvm.amdgcn.raw.buffer.atomic.umax.i32(i32, <4 x i32>, i32, i32, i32 immarg) #0 + +declare i64 @llvm.amdgcn.raw.buffer.atomic.smin.i64(i64, <4 x i32>, i32, i32, i32 immarg) #0 +declare i64 @llvm.amdgcn.raw.buffer.atomic.smax.i64(i64, <4 x i32>, i32, i32, i32 immarg) #0 +declare i64 @llvm.amdgcn.raw.buffer.atomic.umin.i64(i64, <4 x i32>, i32, i32, i32 immarg) #0 +declare i64 @llvm.amdgcn.raw.buffer.atomic.umax.i64(i64, <4 x i32>, i32, i32, i32 immarg) #0 + +attributes #0 = { nounwind } From e6b99fa5b4e7f67930c0aa44b5ea79472fa21187 Mon Sep 17 00:00:00 2001 From: Vy Nguyen Date: Wed, 11 Feb 2026 14:24:42 -0500 Subject: [PATCH 007/705] [LLDB]Fix asan error from pr/179799 (#180996) Fix asan error mentioned in [this comment](https://github.com/llvm/llvm-project/pull/179799#issuecomment-3882710678) This is breaking our internal builds. Details: The field needs to be cleaned up or we'll have a leak. We could also make it a smart pointer if you'd like. But as-is, this is the quickest way to resolve the issue. --- lldb/source/Target/Process.cpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/lldb/source/Target/Process.cpp b/lldb/source/Target/Process.cpp index 67978d3efe723..9371997451beb 100644 --- a/lldb/source/Target/Process.cpp +++ b/lldb/source/Target/Process.cpp @@ -4006,6 +4006,8 @@ void Process::StopPrivateStateThread() { log, "Went to stop the private state thread, but it was already invalid."); } + delete m_current_private_state_thread; + m_current_private_state_thread = nullptr; } void Process::ControlPrivateStateThread(uint32_t signal) { From 2fd5479f2c78780c76424d4d918d69b32bf1573c Mon Sep 17 00:00:00 2001 From: Charles Zablit Date: Wed, 11 Feb 2026 20:26:00 +0100 Subject: [PATCH 008/705] [lldb][windows] fix undeclared identifier error GetModulePath (#180989) `LLDB_PYTHON_DLL_RELATIVE_PATH` can be unset while `LLDB_PYTHON_RUNTIME_LIBRARY_FILENAME` is, since it's set automatically by CMake. This situation causes a build failure which this patch fixes. This issue was introduced in https://github.com/llvm/llvm-project/pull/180784. --- lldb/source/Host/windows/PythonPathSetup/PythonPathSetup.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/lldb/source/Host/windows/PythonPathSetup/PythonPathSetup.cpp b/lldb/source/Host/windows/PythonPathSetup/PythonPathSetup.cpp index 937977a6851c6..b71e10b9c82e7 100644 --- a/lldb/source/Host/windows/PythonPathSetup/PythonPathSetup.cpp +++ b/lldb/source/Host/windows/PythonPathSetup/PythonPathSetup.cpp @@ -19,7 +19,6 @@ using namespace llvm; -#ifdef LLDB_PYTHON_DLL_RELATIVE_PATH static std::string GetModulePath(HMODULE module) { std::vector buffer(MAX_PATH); while (buffer.size() <= PATHCCH_MAX_CCH) { @@ -41,6 +40,7 @@ static std::string GetModulePath(HMODULE module) { /// Returns the full path to the lldb.exe executable. static std::string GetPathToExecutable() { return GetModulePath(NULL); } +#ifdef LLDB_PYTHON_DLL_RELATIVE_PATH bool AddPythonDLLToSearchPath() { std::string path_str = GetPathToExecutable(); if (path_str.empty()) From 5129b3c4494e00fa915a356ec65e649bb7ff6916 Mon Sep 17 00:00:00 2001 From: Maksim Panchenko Date: Wed, 11 Feb 2026 11:35:02 -0800 Subject: [PATCH 009/705] [BOLT] Make FoldedIntoFunction always point to root parent (#180855) After ICF folds functions, FoldedIntoFunction may point to a function that was also folded. Add a post-processing step at the end of ICF to flatten all chains so FoldedIntoFunction always points to the ultimate root parent (a function that is not itself folded). --- bolt/include/bolt/Core/BinaryFunction.h | 3 ++- bolt/lib/Passes/IdenticalCodeFolding.cpp | 13 +++++++++++++ bolt/lib/Rewrite/RewriteInstance.cpp | 11 ++++------- 3 files changed, 19 insertions(+), 8 deletions(-) diff --git a/bolt/include/bolt/Core/BinaryFunction.h b/bolt/include/bolt/Core/BinaryFunction.h index 4637fe311927c..80375b1167b0c 100644 --- a/bolt/include/bolt/Core/BinaryFunction.h +++ b/bolt/include/bolt/Core/BinaryFunction.h @@ -400,7 +400,8 @@ class BinaryFunction { FragmentsSetTy ParentFragments; /// Indicate if the function body was folded into another function. - /// Used by ICF optimization. + /// Used by ICF optimization. Always points to the root parent function + /// (i.e., a function that is not itself folded). BinaryFunction *FoldedIntoFunction{nullptr}; /// All fragments for a parent function. diff --git a/bolt/lib/Passes/IdenticalCodeFolding.cpp b/bolt/lib/Passes/IdenticalCodeFolding.cpp index 1c99a60bfaad3..e412bc1530219 100644 --- a/bolt/lib/Passes/IdenticalCodeFolding.cpp +++ b/bolt/lib/Passes/IdenticalCodeFolding.cpp @@ -605,6 +605,19 @@ Error IdenticalCodeFolding::runOnFunctions(BinaryContext &BC) { } while (NumFoldedLastIteration > 0); + // Flatten folded function chains so FoldedIntoFunction always points + // to the root parent. + for (auto &BFI : BC.getBinaryFunctions()) { + BinaryFunction &BF = BFI.second; + if (!BF.isFolded()) + continue; + BinaryFunction *Parent = BF.getFoldedIntoFunction(); + while (Parent->isFolded()) + Parent = Parent->getFoldedIntoFunction(); + if (Parent != BF.getFoldedIntoFunction()) + BF.setFolded(Parent); + } + LLVM_DEBUG({ // Print functions that are congruent but not identical. for (auto &CBI : CongruentBuckets) { diff --git a/bolt/lib/Rewrite/RewriteInstance.cpp b/bolt/lib/Rewrite/RewriteInstance.cpp index bde326aee2617..4539b7221cb5f 100644 --- a/bolt/lib/Rewrite/RewriteInstance.cpp +++ b/bolt/lib/Rewrite/RewriteInstance.cpp @@ -5162,9 +5162,7 @@ void RewriteInstance::updateELFSymbolTable( auto addExtraSymbols = [&](const BinaryFunction &Function, const ELFSymTy &FunctionSymbol) { if (Function.isFolded()) { - BinaryFunction *ICFParent = Function.getFoldedIntoFunction(); - while (ICFParent->isFolded()) - ICFParent = ICFParent->getFoldedIntoFunction(); + const BinaryFunction *ICFParent = Function.getFoldedIntoFunction(); ELFSymTy ICFSymbol = FunctionSymbol; SmallVector Buf; ICFSymbol.st_name = @@ -5280,8 +5278,7 @@ void RewriteInstance::updateELFSymbolTable( // instead so that the symbol gets updated to the parent's output address. // In non-relocation mode, folded functions are emitted at their original // location, so we keep the original function reference. - // Follow the chain of folded functions to get the final parent. - while (BC->HasRelocations && Function && Function->isFolded()) + if (BC->HasRelocations && Function && Function->isFolded()) Function = Function->getFoldedIntoFunction(); // Ignore false function references, e.g. when the section address matches // the address of the function. @@ -6084,8 +6081,8 @@ uint64_t RewriteInstance::getNewFunctionAddress(uint64_t OldAddress) { return 0; // If this function was folded, its output address is 0 since it wasn't - // emitted. Follow the chain to get the parent function's address. - while (Function->isFolded()) + // emitted. Get the parent function's address. + if (Function->isFolded()) Function = Function->getFoldedIntoFunction(); return Function->getOutputAddress(); From 997f28069d845dcb6d03d8b63cd2173cf818acaa Mon Sep 17 00:00:00 2001 From: Aviral Goel Date: Wed, 11 Feb 2026 11:44:11 -0800 Subject: [PATCH 010/705] [clang][ssaf] Remove VFS support from SerializationFormat (#180891) This change removes support for `VFS` in `SerializationFormat`, introduced in https://github.com/llvm/llvm-project/pull/179516/. `VFS` only provides a virtual input backend and its matched pair, `VirtualOutputBackend`, does not provide support for an `InMemoryOutputBackend`. There was an attempt from our end to implement this support in https://github.com/llvm/llvm-project/pull/179515 but it did not succeed. Supporting virtual reads but not virtual writes makes the SerializationFormat APIs asymmetric. So we have decided to remove `VFS` support altogether. --------- Co-authored-by: Balazs Benics --- .../Serialization/SerializationFormat.h | 5 --- .../SerializationFormatRegistry.h | 11 ++---- .../Serialization/SerializationFormat.cpp | 3 -- .../SerializationFormatRegistry.cpp | 11 ++---- .../Scalable/Registries/FancyAnalysisData.cpp | 1 + .../Registries/MockSerializationFormat.cpp | 11 +++--- .../Registries/MockSerializationFormat.h | 3 +- .../SerializationFormatRegistryTest.cpp | 37 ++++++++++++++----- 8 files changed, 42 insertions(+), 40 deletions(-) diff --git a/clang/include/clang/Analysis/Scalable/Serialization/SerializationFormat.h b/clang/include/clang/Analysis/Scalable/Serialization/SerializationFormat.h index 866fe8bfdcee0..ee220b6be18a6 100644 --- a/clang/include/clang/Analysis/Scalable/Serialization/SerializationFormat.h +++ b/clang/include/clang/Analysis/Scalable/Serialization/SerializationFormat.h @@ -19,7 +19,6 @@ #include "clang/Analysis/Scalable/TUSummary/TUSummary.h" #include "llvm/ADT/StringRef.h" #include "llvm/Support/ExtensibleRTTI.h" -#include "llvm/Support/VirtualFileSystem.h" namespace clang::ssaf { @@ -32,8 +31,6 @@ class EntitySummary; class SerializationFormat : public llvm::RTTIExtends { public: - explicit SerializationFormat( - llvm::IntrusiveRefCntPtr FS); virtual ~SerializationFormat() = default; virtual TUSummary readTUSummary(llvm::StringRef Path) = 0; @@ -50,8 +47,6 @@ class SerializationFormat static const auto &get##FIELD_NAME(const CLASS &X) { return X.FIELD_NAME; } \ static auto &get##FIELD_NAME(CLASS &X) { return X.FIELD_NAME; } #include "clang/Analysis/Scalable/Model/PrivateFieldNames.def" - - llvm::IntrusiveRefCntPtr FS; }; template struct FormatInfoEntry { diff --git a/clang/include/clang/Analysis/Scalable/Serialization/SerializationFormatRegistry.h b/clang/include/clang/Analysis/Scalable/Serialization/SerializationFormatRegistry.h index 40281d549b402..d7e77b9b18f77 100644 --- a/clang/include/clang/Analysis/Scalable/Serialization/SerializationFormatRegistry.h +++ b/clang/include/clang/Analysis/Scalable/Serialization/SerializationFormatRegistry.h @@ -51,21 +51,16 @@ bool isFormatRegistered(llvm::StringRef FormatName); /// This might return null if the construction of the desired /// SerializationFormat failed. /// It's a fatal error if there is no format registered with the name. -std::unique_ptr -makeFormat(llvm::IntrusiveRefCntPtr FS, - llvm::StringRef FormatName); +std::unique_ptr makeFormat(llvm::StringRef FormatName); // Registry for adding new SerializationFormat implementations. -using SerializationFormatRegistry = - llvm::Registry>; +using SerializationFormatRegistry = llvm::Registry; } // namespace clang::ssaf namespace llvm { extern template class CLANG_TEMPLATE_ABI - Registry>; + Registry; } // namespace llvm #endif // CLANG_ANALYSIS_SCALABLE_SERIALIZATION_SERIALIZATION_FORMAT_REGISTRY_H diff --git a/clang/lib/Analysis/Scalable/Serialization/SerializationFormat.cpp b/clang/lib/Analysis/Scalable/Serialization/SerializationFormat.cpp index 75ec263a1051f..07ed1157f87f0 100644 --- a/clang/lib/Analysis/Scalable/Serialization/SerializationFormat.cpp +++ b/clang/lib/Analysis/Scalable/Serialization/SerializationFormat.cpp @@ -11,6 +11,3 @@ using namespace clang::ssaf; char SerializationFormat::ID = 0; -SerializationFormat::SerializationFormat( - llvm::IntrusiveRefCntPtr FS) - : FS(FS) {} diff --git a/clang/lib/Analysis/Scalable/Serialization/SerializationFormatRegistry.cpp b/clang/lib/Analysis/Scalable/Serialization/SerializationFormatRegistry.cpp index f31a154866238..7cfa4370aa9ed 100644 --- a/clang/lib/Analysis/Scalable/Serialization/SerializationFormatRegistry.cpp +++ b/clang/lib/Analysis/Scalable/Serialization/SerializationFormatRegistry.cpp @@ -12,11 +12,7 @@ using namespace clang; using namespace ssaf; -// FIXME: LLVM_INSTANTIATE_REGISTRY can't be used here because it drops extra -// type parameters. -template class CLANG_EXPORT_TEMPLATE - llvm::Registry>; +LLVM_INSTANTIATE_REGISTRY(SerializationFormatRegistry) bool ssaf::isFormatRegistered(llvm::StringRef FormatName) { for (const auto &Entry : SerializationFormatRegistry::entries()) @@ -26,11 +22,10 @@ bool ssaf::isFormatRegistered(llvm::StringRef FormatName) { } std::unique_ptr -ssaf::makeFormat(llvm::IntrusiveRefCntPtr FS, - llvm::StringRef FormatName) { +ssaf::makeFormat(llvm::StringRef FormatName) { for (const auto &Entry : SerializationFormatRegistry::entries()) if (Entry.getName() == FormatName) - return Entry.instantiate(std::move(FS)); + return Entry.instantiate(); assert(false && "Unknown SerializationFormat name"); return nullptr; } diff --git a/clang/unittests/Analysis/Scalable/Registries/FancyAnalysisData.cpp b/clang/unittests/Analysis/Scalable/Registries/FancyAnalysisData.cpp index 1198e7a927a28..ac4e625aaf671 100644 --- a/clang/unittests/Analysis/Scalable/Registries/FancyAnalysisData.cpp +++ b/clang/unittests/Analysis/Scalable/Registries/FancyAnalysisData.cpp @@ -8,6 +8,7 @@ #include "Registries/MockSerializationFormat.h" #include "clang/Analysis/Scalable/TUSummary/EntitySummary.h" +#include "llvm/Support/Casting.h" #include "llvm/Support/Registry.h" using namespace clang; diff --git a/clang/unittests/Analysis/Scalable/Registries/MockSerializationFormat.cpp b/clang/unittests/Analysis/Scalable/Registries/MockSerializationFormat.cpp index 127e61b26c0cb..719db61486df5 100644 --- a/clang/unittests/Analysis/Scalable/Registries/MockSerializationFormat.cpp +++ b/clang/unittests/Analysis/Scalable/Registries/MockSerializationFormat.cpp @@ -17,6 +17,7 @@ #include "llvm/ADT/SmallVector.h" #include "llvm/ADT/StringRef.h" #include "llvm/Support/FileSystem.h" +#include "llvm/Support/MemoryBuffer.h" #include "llvm/Support/raw_ostream.h" #include #include @@ -28,9 +29,8 @@ using namespace ssaf; char MockSerializationFormat::ID = 0; -MockSerializationFormat::MockSerializationFormat( - llvm::IntrusiveRefCntPtr FS) - : llvm::RTTIExtends(FS) { +MockSerializationFormat::MockSerializationFormat() + : llvm::RTTIExtends() { for (const auto &FormatInfoEntry : llvm::Registry::entries()) { std::unique_ptr Info = FormatInfoEntry.instantiate(); bool Inserted = FormatInfos.try_emplace(Info->ForSummary, *Info).second; @@ -46,7 +46,7 @@ TUSummary MockSerializationFormat::readTUSummary(llvm::StringRef Path) { BuildNamespace NS(BuildNamespaceKind::CompilationUnit, "Mock.cpp"); TUSummary Summary(NS); - auto ManifestFile = FS->getBufferForFile(Path + "/analyses.txt"); + auto ManifestFile = llvm::MemoryBuffer::getFile(Path + "/analyses.txt"); assert(ManifestFile); // TODO Handle error. llvm::StringRef ManifestFileContent = (*ManifestFile)->getBuffer(); @@ -56,7 +56,8 @@ TUSummary MockSerializationFormat::readTUSummary(llvm::StringRef Path) { for (llvm::StringRef Analysis : Analyses) { SummaryName Name(Analysis.str()); - auto InputFile = FS->getBufferForFile(Path + "/" + Name.str() + ".special"); + auto InputFile = + llvm::MemoryBuffer::getFile(Path + "/" + Name.str() + ".special"); assert(InputFile); auto InfoIt = FormatInfos.find(Name); if (InfoIt == FormatInfos.end()) { diff --git a/clang/unittests/Analysis/Scalable/Registries/MockSerializationFormat.h b/clang/unittests/Analysis/Scalable/Registries/MockSerializationFormat.h index 25fb2f3b20551..8c287c2143621 100644 --- a/clang/unittests/Analysis/Scalable/Registries/MockSerializationFormat.h +++ b/clang/unittests/Analysis/Scalable/Registries/MockSerializationFormat.h @@ -20,8 +20,7 @@ namespace clang::ssaf { class MockSerializationFormat final : public llvm::RTTIExtends { public: - explicit MockSerializationFormat( - llvm::IntrusiveRefCntPtr FS); + MockSerializationFormat(); TUSummary readTUSummary(llvm::StringRef Path) override; diff --git a/clang/unittests/Analysis/Scalable/Registries/SerializationFormatRegistryTest.cpp b/clang/unittests/Analysis/Scalable/Registries/SerializationFormatRegistryTest.cpp index 484c38309e3d8..7f17cdda5a082 100644 --- a/clang/unittests/Analysis/Scalable/Registries/SerializationFormatRegistryTest.cpp +++ b/clang/unittests/Analysis/Scalable/Registries/SerializationFormatRegistryTest.cpp @@ -13,7 +13,7 @@ #include "llvm/Support/FileSystem.h" #include "llvm/Support/MemoryBuffer.h" #include "llvm/Support/Path.h" -#include "llvm/Support/VirtualFileSystem.h" +#include "llvm/Support/raw_ostream.h" #include "gtest/gtest.h" #include @@ -58,21 +58,40 @@ TEST(SerializationFormatRegistryTest, EnumeratingRegistryEntries) { } TEST(SerializationFormatRegistryTest, Roundtrip) { - auto Inputs = makeIntrusiveRefCnt(); - Inputs->addFile("input/analyses.txt", /*ModificationTime=*/{}, - MemoryBuffer::getMemBufferCopy("FancyAnalysis\n")); - Inputs->addFile("input/FancyAnalysis.special", /*ModificationTime=*/{}, - MemoryBuffer::getMemBufferCopy("Some FancyAnalysisData...")); + // Create temporary input directory + SmallString<128> InputDir; + std::error_code EC = sys::fs::createUniqueDirectory("ssaf-input", InputDir); + ASSERT_FALSE(EC) << "Failed to create input directory: " << EC.message(); + llvm::scope_exit CleanupInputOnExit( + [&] { sys::fs::remove_directories(InputDir); }); + + // Create input files + SmallString<128> AnalysesFile = InputDir; + sys::path::append(AnalysesFile, "analyses.txt"); + { + raw_fd_ostream OS(AnalysesFile, EC); + ASSERT_FALSE(EC) << "Failed to create analyses.txt: " << EC.message(); + OS << "FancyAnalysis\n"; + } + + SmallString<128> FancyAnalysisFile = InputDir; + sys::path::append(FancyAnalysisFile, "FancyAnalysis.special"); + { + raw_fd_ostream OS(FancyAnalysisFile, EC); + ASSERT_FALSE(EC) << "Failed to create FancyAnalysis.special: " + << EC.message(); + OS << "Some FancyAnalysisData..."; + } std::unique_ptr Format = - makeFormat(Inputs, "MockSerializationFormat"); + makeFormat("MockSerializationFormat"); ASSERT_TRUE(Format); - TUSummary LoadedSummary = Format->readTUSummary("input"); + TUSummary LoadedSummary = Format->readTUSummary(InputDir); // Create a temporary output directory SmallString<128> OutputDir; - std::error_code EC = sys::fs::createUniqueDirectory("ssaf-test", OutputDir); + EC = sys::fs::createUniqueDirectory("ssaf-test", OutputDir); ASSERT_FALSE(EC) << "Failed to create temporary directory: " << EC.message(); llvm::scope_exit CleanupOnExit( [&] { sys::fs::remove_directories(OutputDir); }); From 6fdf93c5d812197dd5dc9362f324b944105b2e8e Mon Sep 17 00:00:00 2001 From: Nikita Kornev Date: Wed, 11 Feb 2026 20:49:12 +0100 Subject: [PATCH 011/705] Update requirements_formatting.txt.in (#180963) At https://github.com/intel/llvm dependabot reports a few alerts regarding use of: - urllib3 < v2.6.3 (CVE-2026-21441) - pynacl < 1.6.2 (CVE-2025-69277) - cryptography <= 46.0.4 (CVE-2026-26007) Update requirements_formatting.txt.in & re-generate requirements_formatting.txt. --- llvm/utils/git/requirements_formatting.txt | 347 ++++++++++++------ llvm/utils/git/requirements_formatting.txt.in | 3 + 2 files changed, 232 insertions(+), 118 deletions(-) diff --git a/llvm/utils/git/requirements_formatting.txt b/llvm/utils/git/requirements_formatting.txt index 18e2626c79460..c1bed42e5f8c0 100644 --- a/llvm/utils/git/requirements_formatting.txt +++ b/llvm/utils/git/requirements_formatting.txt @@ -1,5 +1,5 @@ # -# This file is autogenerated by pip-compile with Python 3.11 +# This file is autogenerated by pip-compile with Python 3.10 # by the following command: # # pip-compile --generate-hashes --output-file=requirements_formatting.txt requirements_formatting.txt.in @@ -34,74 +34,91 @@ certifi==2024.7.4 \ --hash=sha256:5a1e7645bc0ec61a09e26c36f6106dd4cf40c6db3a1fb6352b0244e7fb057c7b \ --hash=sha256:c198e21b1289c2ab85ee4e67bb4b4ef3ead0892059901a8d5b622f24a1101e90 # via requests -cffi==1.17.0 \ - --hash=sha256:011aff3524d578a9412c8b3cfaa50f2c0bd78e03eb7af7aa5e0df59b158efb2f \ - --hash=sha256:0a048d4f6630113e54bb4b77e315e1ba32a5a31512c31a273807d0027a7e69ab \ - --hash=sha256:0bb15e7acf8ab35ca8b24b90af52c8b391690ef5c4aec3d31f38f0d37d2cc499 \ - --hash=sha256:0d46ee4764b88b91f16661a8befc6bfb24806d885e27436fdc292ed7e6f6d058 \ - --hash=sha256:0e60821d312f99d3e1569202518dddf10ae547e799d75aef3bca3a2d9e8ee693 \ - --hash=sha256:0fdacad9e0d9fc23e519efd5ea24a70348305e8d7d85ecbb1a5fa66dc834e7fb \ - --hash=sha256:14b9cbc8f7ac98a739558eb86fabc283d4d564dafed50216e7f7ee62d0d25377 \ - --hash=sha256:17c6d6d3260c7f2d94f657e6872591fe8733872a86ed1345bda872cfc8c74885 \ - --hash=sha256:1a2ddbac59dc3716bc79f27906c010406155031a1c801410f1bafff17ea304d2 \ - --hash=sha256:2404f3de742f47cb62d023f0ba7c5a916c9c653d5b368cc966382ae4e57da401 \ - --hash=sha256:24658baf6224d8f280e827f0a50c46ad819ec8ba380a42448e24459daf809cf4 \ - --hash=sha256:24aa705a5f5bd3a8bcfa4d123f03413de5d86e497435693b638cbffb7d5d8a1b \ - --hash=sha256:2770bb0d5e3cc0e31e7318db06efcbcdb7b31bcb1a70086d3177692a02256f59 \ - --hash=sha256:331ad15c39c9fe9186ceaf87203a9ecf5ae0ba2538c9e898e3a6967e8ad3db6f \ - --hash=sha256:3aa9d43b02a0c681f0bfbc12d476d47b2b2b6a3f9287f11ee42989a268a1833c \ - --hash=sha256:41f4915e09218744d8bae14759f983e466ab69b178de38066f7579892ff2a555 \ - --hash=sha256:4304d4416ff032ed50ad6bb87416d802e67139e31c0bde4628f36a47a3164bfa \ - --hash=sha256:435a22d00ec7d7ea533db494da8581b05977f9c37338c80bc86314bec2619424 \ - --hash=sha256:45f7cd36186db767d803b1473b3c659d57a23b5fa491ad83c6d40f2af58e4dbb \ - --hash=sha256:48b389b1fd5144603d61d752afd7167dfd205973a43151ae5045b35793232aa2 \ - --hash=sha256:4e67d26532bfd8b7f7c05d5a766d6f437b362c1bf203a3a5ce3593a645e870b8 \ - --hash=sha256:516a405f174fd3b88829eabfe4bb296ac602d6a0f68e0d64d5ac9456194a5b7e \ - --hash=sha256:5ba5c243f4004c750836f81606a9fcb7841f8874ad8f3bf204ff5e56332b72b9 \ - --hash=sha256:5bdc0f1f610d067c70aa3737ed06e2726fd9d6f7bfee4a351f4c40b6831f4e82 \ - --hash=sha256:6107e445faf057c118d5050560695e46d272e5301feffda3c41849641222a828 \ - --hash=sha256:6327b572f5770293fc062a7ec04160e89741e8552bf1c358d1a23eba68166759 \ - --hash=sha256:669b29a9eca6146465cc574659058ed949748f0809a2582d1f1a324eb91054dc \ - --hash=sha256:6ce01337d23884b21c03869d2f68c5523d43174d4fc405490eb0091057943118 \ - --hash=sha256:6d872186c1617d143969defeadac5a904e6e374183e07977eedef9c07c8953bf \ - --hash=sha256:6f76a90c345796c01d85e6332e81cab6d70de83b829cf1d9762d0a3da59c7932 \ - --hash=sha256:70d2aa9fb00cf52034feac4b913181a6e10356019b18ef89bc7c12a283bf5f5a \ - --hash=sha256:7cbc78dc018596315d4e7841c8c3a7ae31cc4d638c9b627f87d52e8abaaf2d29 \ - --hash=sha256:856bf0924d24e7f93b8aee12a3a1095c34085600aa805693fb7f5d1962393206 \ - --hash=sha256:8a98748ed1a1df4ee1d6f927e151ed6c1a09d5ec21684de879c7ea6aa96f58f2 \ - --hash=sha256:93a7350f6706b31f457c1457d3a3259ff9071a66f312ae64dc024f049055f72c \ - --hash=sha256:964823b2fc77b55355999ade496c54dde161c621cb1f6eac61dc30ed1b63cd4c \ - --hash=sha256:a003ac9edc22d99ae1286b0875c460351f4e101f8c9d9d2576e78d7e048f64e0 \ - --hash=sha256:a0ce71725cacc9ebf839630772b07eeec220cbb5f03be1399e0457a1464f8e1a \ - --hash=sha256:a47eef975d2b8b721775a0fa286f50eab535b9d56c70a6e62842134cf7841195 \ - --hash=sha256:a8b5b9712783415695663bd463990e2f00c6750562e6ad1d28e072a611c5f2a6 \ - --hash=sha256:a9015f5b8af1bb6837a3fcb0cdf3b874fe3385ff6274e8b7925d81ccaec3c5c9 \ - --hash=sha256:aec510255ce690d240f7cb23d7114f6b351c733a74c279a84def763660a2c3bc \ - --hash=sha256:b00e7bcd71caa0282cbe3c90966f738e2db91e64092a877c3ff7f19a1628fdcb \ - --hash=sha256:b50aaac7d05c2c26dfd50c3321199f019ba76bb650e346a6ef3616306eed67b0 \ - --hash=sha256:b7b6ea9e36d32582cda3465f54c4b454f62f23cb083ebc7a94e2ca6ef011c3a7 \ - --hash=sha256:bb9333f58fc3a2296fb1d54576138d4cf5d496a2cc118422bd77835e6ae0b9cb \ - --hash=sha256:c1c13185b90bbd3f8b5963cd8ce7ad4ff441924c31e23c975cb150e27c2bf67a \ - --hash=sha256:c3b8bd3133cd50f6b637bb4322822c94c5ce4bf0d724ed5ae70afce62187c492 \ - --hash=sha256:c5d97162c196ce54af6700949ddf9409e9833ef1003b4741c2b39ef46f1d9720 \ - --hash=sha256:c815270206f983309915a6844fe994b2fa47e5d05c4c4cef267c3b30e34dbe42 \ - --hash=sha256:cab2eba3830bf4f6d91e2d6718e0e1c14a2f5ad1af68a89d24ace0c6b17cced7 \ - --hash=sha256:d1df34588123fcc88c872f5acb6f74ae59e9d182a2707097f9e28275ec26a12d \ - --hash=sha256:d6bdcd415ba87846fd317bee0774e412e8792832e7805938987e4ede1d13046d \ - --hash=sha256:db9a30ec064129d605d0f1aedc93e00894b9334ec74ba9c6bdd08147434b33eb \ - --hash=sha256:dbc183e7bef690c9abe5ea67b7b60fdbca81aa8da43468287dae7b5c046107d4 \ - --hash=sha256:dca802c8db0720ce1c49cce1149ff7b06e91ba15fa84b1d59144fef1a1bc7ac2 \ - --hash=sha256:dec6b307ce928e8e112a6bb9921a1cb00a0e14979bf28b98e084a4b8a742bd9b \ - --hash=sha256:df8bb0010fdd0a743b7542589223a2816bdde4d94bb5ad67884348fa2c1c67e8 \ - --hash=sha256:e4094c7b464cf0a858e75cd14b03509e84789abf7b79f8537e6a72152109c76e \ - --hash=sha256:e4760a68cab57bfaa628938e9c2971137e05ce48e762a9cb53b76c9b569f1204 \ - --hash=sha256:eb09b82377233b902d4c3fbeeb7ad731cdab579c6c6fda1f763cd779139e47c3 \ - --hash=sha256:eb862356ee9391dc5a0b3cbc00f416b48c1b9a52d252d898e5b7696a5f9fe150 \ - --hash=sha256:ef9528915df81b8f4c7612b19b8628214c65c9b7f74db2e34a646a0a2a0da2d4 \ - --hash=sha256:f3157624b7558b914cb039fd1af735e5e8049a87c817cc215109ad1c8779df76 \ - --hash=sha256:f3e0992f23bbb0be00a921eae5363329253c3b86287db27092461c887b791e5e \ - --hash=sha256:f9338cc05451f1942d0d8203ec2c346c830f8e86469903d5126c1f0a13a2bcbb \ - --hash=sha256:ffef8fd58a36fb5f1196919638f73dd3ae0db1a878982b27a9a5a176ede4ba91 +cffi==2.0.0 \ + --hash=sha256:00bdf7acc5f795150faa6957054fbbca2439db2f775ce831222b66f192f03beb \ + --hash=sha256:07b271772c100085dd28b74fa0cd81c8fb1a3ba18b21e03d7c27f3436a10606b \ + --hash=sha256:087067fa8953339c723661eda6b54bc98c5625757ea62e95eb4898ad5e776e9f \ + --hash=sha256:0a1527a803f0a659de1af2e1fd700213caba79377e27e4693648c2923da066f9 \ + --hash=sha256:0cf2d91ecc3fcc0625c2c530fe004f82c110405f101548512cce44322fa8ac44 \ + --hash=sha256:0f6084a0ea23d05d20c3edcda20c3d006f9b6f3fefeac38f59262e10cef47ee2 \ + --hash=sha256:12873ca6cb9b0f0d3a0da705d6086fe911591737a59f28b7936bdfed27c0d47c \ + --hash=sha256:19f705ada2530c1167abacb171925dd886168931e0a7b78f5bffcae5c6b5be75 \ + --hash=sha256:1cd13c99ce269b3ed80b417dcd591415d3372bcac067009b6e0f59c7d4015e65 \ + --hash=sha256:1e3a615586f05fc4065a8b22b8152f0c1b00cdbc60596d187c2a74f9e3036e4e \ + --hash=sha256:1f72fb8906754ac8a2cc3f9f5aaa298070652a0ffae577e0ea9bd480dc3c931a \ + --hash=sha256:1fc9ea04857caf665289b7a75923f2c6ed559b8298a1b8c49e59f7dd95c8481e \ + --hash=sha256:203a48d1fb583fc7d78a4c6655692963b860a417c0528492a6bc21f1aaefab25 \ + --hash=sha256:2081580ebb843f759b9f617314a24ed5738c51d2aee65d31e02f6f7a2b97707a \ + --hash=sha256:21d1152871b019407d8ac3985f6775c079416c282e431a4da6afe7aefd2bccbe \ + --hash=sha256:24b6f81f1983e6df8db3adc38562c83f7d4a0c36162885ec7f7b77c7dcbec97b \ + --hash=sha256:256f80b80ca3853f90c21b23ee78cd008713787b1b1e93eae9f3d6a7134abd91 \ + --hash=sha256:28a3a209b96630bca57cce802da70c266eb08c6e97e5afd61a75611ee6c64592 \ + --hash=sha256:2c8f814d84194c9ea681642fd164267891702542f028a15fc97d4674b6206187 \ + --hash=sha256:2de9a304e27f7596cd03d16f1b7c72219bd944e99cc52b84d0145aefb07cbd3c \ + --hash=sha256:38100abb9d1b1435bc4cc340bb4489635dc2f0da7456590877030c9b3d40b0c1 \ + --hash=sha256:3925dd22fa2b7699ed2617149842d2e6adde22b262fcbfada50e3d195e4b3a94 \ + --hash=sha256:3e17ed538242334bf70832644a32a7aae3d83b57567f9fd60a26257e992b79ba \ + --hash=sha256:3e837e369566884707ddaf85fc1744b47575005c0a229de3327f8f9a20f4efeb \ + --hash=sha256:3f4d46d8b35698056ec29bca21546e1551a205058ae1a181d871e278b0b28165 \ + --hash=sha256:44d1b5909021139fe36001ae048dbdde8214afa20200eda0f64c068cac5d5529 \ + --hash=sha256:45d5e886156860dc35862657e1494b9bae8dfa63bf56796f2fb56e1679fc0bca \ + --hash=sha256:4647afc2f90d1ddd33441e5b0e85b16b12ddec4fca55f0d9671fef036ecca27c \ + --hash=sha256:4671d9dd5ec934cb9a73e7ee9676f9362aba54f7f34910956b84d727b0d73fb6 \ + --hash=sha256:53f77cbe57044e88bbd5ed26ac1d0514d2acf0591dd6bb02a3ae37f76811b80c \ + --hash=sha256:5eda85d6d1879e692d546a078b44251cdd08dd1cfb98dfb77b670c97cee49ea0 \ + --hash=sha256:5fed36fccc0612a53f1d4d9a816b50a36702c28a2aa880cb8a122b3466638743 \ + --hash=sha256:61d028e90346df14fedc3d1e5441df818d095f3b87d286825dfcbd6459b7ef63 \ + --hash=sha256:66f011380d0e49ed280c789fbd08ff0d40968ee7b665575489afa95c98196ab5 \ + --hash=sha256:6824f87845e3396029f3820c206e459ccc91760e8fa24422f8b0c3d1731cbec5 \ + --hash=sha256:6c6c373cfc5c83a975506110d17457138c8c63016b563cc9ed6e056a82f13ce4 \ + --hash=sha256:6d02d6655b0e54f54c4ef0b94eb6be0607b70853c45ce98bd278dc7de718be5d \ + --hash=sha256:6d50360be4546678fc1b79ffe7a66265e28667840010348dd69a314145807a1b \ + --hash=sha256:730cacb21e1bdff3ce90babf007d0a0917cc3e6492f336c2f0134101e0944f93 \ + --hash=sha256:737fe7d37e1a1bffe70bd5754ea763a62a066dc5913ca57e957824b72a85e205 \ + --hash=sha256:74a03b9698e198d47562765773b4a8309919089150a0bb17d829ad7b44b60d27 \ + --hash=sha256:7553fb2090d71822f02c629afe6042c299edf91ba1bf94951165613553984512 \ + --hash=sha256:7a66c7204d8869299919db4d5069a82f1561581af12b11b3c9f48c584eb8743d \ + --hash=sha256:7cc09976e8b56f8cebd752f7113ad07752461f48a58cbba644139015ac24954c \ + --hash=sha256:81afed14892743bbe14dacb9e36d9e0e504cd204e0b165062c488942b9718037 \ + --hash=sha256:8941aaadaf67246224cee8c3803777eed332a19d909b47e29c9842ef1e79ac26 \ + --hash=sha256:89472c9762729b5ae1ad974b777416bfda4ac5642423fa93bd57a09204712322 \ + --hash=sha256:8ea985900c5c95ce9db1745f7933eeef5d314f0565b27625d9a10ec9881e1bfb \ + --hash=sha256:8eca2a813c1cb7ad4fb74d368c2ffbbb4789d377ee5bb8df98373c2cc0dee76c \ + --hash=sha256:92b68146a71df78564e4ef48af17551a5ddd142e5190cdf2c5624d0c3ff5b2e8 \ + --hash=sha256:9332088d75dc3241c702d852d4671613136d90fa6881da7d770a483fd05248b4 \ + --hash=sha256:94698a9c5f91f9d138526b48fe26a199609544591f859c870d477351dc7b2414 \ + --hash=sha256:9a67fc9e8eb39039280526379fb3a70023d77caec1852002b4da7e8b270c4dd9 \ + --hash=sha256:9de40a7b0323d889cf8d23d1ef214f565ab154443c42737dfe52ff82cf857664 \ + --hash=sha256:a05d0c237b3349096d3981b727493e22147f934b20f6f125a3eba8f994bec4a9 \ + --hash=sha256:afb8db5439b81cf9c9d0c80404b60c3cc9c3add93e114dcae767f1477cb53775 \ + --hash=sha256:b18a3ed7d5b3bd8d9ef7a8cb226502c6bf8308df1525e1cc676c3680e7176739 \ + --hash=sha256:b1e74d11748e7e98e2f426ab176d4ed720a64412b6a15054378afdb71e0f37dc \ + --hash=sha256:b21e08af67b8a103c71a250401c78d5e0893beff75e28c53c98f4de42f774062 \ + --hash=sha256:b4c854ef3adc177950a8dfc81a86f5115d2abd545751a304c5bcf2c2c7283cfe \ + --hash=sha256:b882b3df248017dba09d6b16defe9b5c407fe32fc7c65a9c69798e6175601be9 \ + --hash=sha256:baf5215e0ab74c16e2dd324e8ec067ef59e41125d3eade2b863d294fd5035c92 \ + --hash=sha256:c649e3a33450ec82378822b3dad03cc228b8f5963c0c12fc3b1e0ab940f768a5 \ + --hash=sha256:c654de545946e0db659b3400168c9ad31b5d29593291482c43e3564effbcee13 \ + --hash=sha256:c6638687455baf640e37344fe26d37c404db8b80d037c3d29f58fe8d1c3b194d \ + --hash=sha256:c8d3b5532fc71b7a77c09192b4a5a200ea992702734a2e9279a37f2478236f26 \ + --hash=sha256:cb527a79772e5ef98fb1d700678fe031e353e765d1ca2d409c92263c6d43e09f \ + --hash=sha256:cf364028c016c03078a23b503f02058f1814320a56ad535686f90565636a9495 \ + --hash=sha256:d48a880098c96020b02d5a1f7d9251308510ce8858940e6fa99ece33f610838b \ + --hash=sha256:d68b6cef7827e8641e8ef16f4494edda8b36104d79773a334beaa1e3521430f6 \ + --hash=sha256:d9b29c1f0ae438d5ee9acb31cadee00a58c46cc9c0b2f9038c6b0b3470877a8c \ + --hash=sha256:d9b97165e8aed9272a6bb17c01e3cc5871a594a446ebedc996e2397a1c1ea8ef \ + --hash=sha256:da68248800ad6320861f129cd9c1bf96ca849a2771a59e0344e88681905916f5 \ + --hash=sha256:da902562c3e9c550df360bfa53c035b2f241fed6d9aef119048073680ace4a18 \ + --hash=sha256:dbd5c7a25a7cb98f5ca55d258b103a2054f859a46ae11aaf23134f9cc0d356ad \ + --hash=sha256:dd4f05f54a52fb558f1ba9f528228066954fee3ebe629fc1660d874d040ae5a3 \ + --hash=sha256:de8dad4425a6ca6e4e5e297b27b5c824ecc7581910bf9aee86cb6835e6812aa7 \ + --hash=sha256:e11e82b744887154b182fd3e7e8512418446501191994dbf9c9fc1f32cc8efd5 \ + --hash=sha256:e6e73b9e02893c764e7e8d5bb5ce277f1a009cd5243f8228f75f842bf937c534 \ + --hash=sha256:f73b96c41e3b2adedc34a7356e64c8eb96e03a3782b535e043a986276ce12a49 \ + --hash=sha256:f93fd8e5c8c0a4aa1f424d6173f14a892044054871c771f8566e4008eaa359d2 \ + --hash=sha256:fc33c5141b55ed366cfaad382df24fe7dcbc686de5be719b207bb248e3053dc5 \ + --hash=sha256:fc7de24befaeae77ba923797c7c87834c73648a05a4bde34b3b7e5588973a453 \ + --hash=sha256:fe562eb1a64e67dd297ccc4f5addea2501664954f2692b69a76449ec7913ecbf # via # cryptography # pynacl @@ -201,39 +218,59 @@ click==8.1.7 \ --hash=sha256:ae74fb96c20a0277a1d615f1e4d73c8414f5a98db8b799a7931d1582f3390c28 \ --hash=sha256:ca9853ad459e787e2192211578cc907e7594e294c7ccc834310722b41b9ca6de # via black -colorama==0.4.6 \ - --hash=sha256:08695f5cb7ed6e0531a20572697297273c47b8cae5a63ffc6d6ed5c201be6e44 \ - --hash=sha256:4f1d9991f5acc0ca119f9d443620b77f9d6b33703e51011c16baf57afb285fc6 - # via click -cryptography==43.0.0 \ - --hash=sha256:0663585d02f76929792470451a5ba64424acc3cd5227b03921dab0e2f27b1709 \ - --hash=sha256:08a24a7070b2b6804c1940ff0f910ff728932a9d0e80e7814234269f9d46d069 \ - --hash=sha256:232ce02943a579095a339ac4b390fbbe97f5b5d5d107f8a08260ea2768be8cc2 \ - --hash=sha256:2905ccf93a8a2a416f3ec01b1a7911c3fe4073ef35640e7ee5296754e30b762b \ - --hash=sha256:299d3da8e00b7e2b54bb02ef58d73cd5f55fb31f33ebbf33bd00d9aa6807df7e \ - --hash=sha256:2c6d112bf61c5ef44042c253e4859b3cbbb50df2f78fa8fae6747a7814484a70 \ - --hash=sha256:31e44a986ceccec3d0498e16f3d27b2ee5fdf69ce2ab89b52eaad1d2f33d8778 \ - --hash=sha256:3d9a1eca329405219b605fac09ecfc09ac09e595d6def650a437523fcd08dd22 \ - --hash=sha256:3dcdedae5c7710b9f97ac6bba7e1052b95c7083c9d0e9df96e02a1932e777895 \ - --hash=sha256:47ca71115e545954e6c1d207dd13461ab81f4eccfcb1345eac874828b5e3eaaf \ - --hash=sha256:4a997df8c1c2aae1e1e5ac49c2e4f610ad037fc5a3aadc7b64e39dea42249431 \ - --hash=sha256:51956cf8730665e2bdf8ddb8da0056f699c1a5715648c1b0144670c1ba00b48f \ - --hash=sha256:5bcb8a5620008a8034d39bce21dc3e23735dfdb6a33a06974739bfa04f853947 \ - --hash=sha256:64c3f16e2a4fc51c0d06af28441881f98c5d91009b8caaff40cf3548089e9c74 \ - --hash=sha256:6e2b11c55d260d03a8cf29ac9b5e0608d35f08077d8c087be96287f43af3ccdc \ - --hash=sha256:7b3f5fe74a5ca32d4d0f302ffe6680fcc5c28f8ef0dc0ae8f40c0f3a1b4fca66 \ - --hash=sha256:844b6d608374e7d08f4f6e6f9f7b951f9256db41421917dfb2d003dde4cd6b66 \ - --hash=sha256:9a8d6802e0825767476f62aafed40532bd435e8a5f7d23bd8b4f5fd04cc80ecf \ - --hash=sha256:aae4d918f6b180a8ab8bf6511a419473d107df4dbb4225c7b48c5c9602c38c7f \ - --hash=sha256:ac1955ce000cb29ab40def14fd1bbfa7af2017cca696ee696925615cafd0dce5 \ - --hash=sha256:b88075ada2d51aa9f18283532c9f60e72170041bba88d7f37e49cbb10275299e \ - --hash=sha256:cb013933d4c127349b3948aa8aaf2f12c0353ad0eccd715ca789c8a0f671646f \ - --hash=sha256:cc70b4b581f28d0a254d006f26949245e3657d40d8857066c2ae22a61222ef55 \ - --hash=sha256:e9c5266c432a1e23738d178e51c2c7a5e2ddf790f248be939448c0ba2021f9d1 \ - --hash=sha256:ea9e57f8ea880eeea38ab5abf9fbe39f923544d7884228ec67d666abd60f5a47 \ - --hash=sha256:ee0c405832ade84d4de74b9029bedb7b31200600fa524d218fc29bfa371e97f5 \ - --hash=sha256:fdcb265de28585de5b859ae13e3846a8e805268a823a12a4da2597f1f5afc9f0 - # via pyjwt +cryptography==46.0.5 \ + --hash=sha256:02f547fce831f5096c9a567fd41bc12ca8f11df260959ecc7c3202555cc47a72 \ + --hash=sha256:039917b0dc418bb9f6edce8a906572d69e74bd330b0b3fea4f79dab7f8ddd235 \ + --hash=sha256:1abfdb89b41c3be0365328a410baa9df3ff8a9110fb75e7b52e66803ddabc9a9 \ + --hash=sha256:2ae6971afd6246710480e3f15824ed3029a60fc16991db250034efd0b9fb4356 \ + --hash=sha256:2b7a67c9cd56372f3249b39699f2ad479f6991e62ea15800973b956f4b73e257 \ + --hash=sha256:351695ada9ea9618b3500b490ad54c739860883df6c1f555e088eaf25b1bbaad \ + --hash=sha256:38946c54b16c885c72c4f59846be9743d699eee2b69b6988e0a00a01f46a61a4 \ + --hash=sha256:3b4995dc971c9fb83c25aa44cf45f02ba86f71ee600d81091c2f0cbae116b06c \ + --hash=sha256:3ce58ba46e1bc2aac4f7d9290223cead56743fa6ab94a5d53292ffaac6a91614 \ + --hash=sha256:3ee190460e2fbe447175cda91b88b84ae8322a104fc27766ad09428754a618ed \ + --hash=sha256:4108d4c09fbbf2789d0c926eb4152ae1760d5a2d97612b92d508d96c861e4d31 \ + --hash=sha256:420d0e909050490d04359e7fdb5ed7e667ca5c3c402b809ae2563d7e66a92229 \ + --hash=sha256:47fb8a66058b80e509c47118ef8a75d14c455e81ac369050f20ba0d23e77fee0 \ + --hash=sha256:4c3341037c136030cb46e4b1e17b7418ea4cbd9dd207e4a6f3b2b24e0d4ac731 \ + --hash=sha256:4d7e3d356b8cd4ea5aff04f129d5f66ebdc7b6f8eae802b93739ed520c47c79b \ + --hash=sha256:4d8ae8659ab18c65ced284993c2265910f6c9e650189d4e3f68445ef82a810e4 \ + --hash=sha256:4e817a8920bfbcff8940ecfd60f23d01836408242b30f1a708d93198393a80b4 \ + --hash=sha256:50bfb6925eff619c9c023b967d5b77a54e04256c4281b0e21336a130cd7fc263 \ + --hash=sha256:556e106ee01aa13484ce9b0239bca667be5004efb0aabbed28d353df86445595 \ + --hash=sha256:582f5fcd2afa31622f317f80426a027f30dc792e9c80ffee87b993200ea115f1 \ + --hash=sha256:5be7bf2fb40769e05739dd0046e7b26f9d4670badc7b032d6ce4db64dddc0678 \ + --hash=sha256:60ee7e19e95104d4c03871d7d7dfb3d22ef8a9b9c6778c94e1c8fcc8365afd48 \ + --hash=sha256:61aa400dce22cb001a98014f647dc21cda08f7915ceb95df0c9eaf84b4b6af76 \ + --hash=sha256:68f68d13f2e1cb95163fa3b4db4bf9a159a418f5f6e7242564fc75fcae667fd0 \ + --hash=sha256:7d1f30a86d2757199cb2d56e48cce14deddf1f9c95f1ef1b64ee91ea43fe2e18 \ + --hash=sha256:7d731d4b107030987fd61a7f8ab512b25b53cef8f233a97379ede116f30eb67d \ + --hash=sha256:803812e111e75d1aa73690d2facc295eaefd4439be1023fefc4995eaea2af90d \ + --hash=sha256:80a8d7bfdf38f87ca30a5391c0c9ce4ed2926918e017c29ddf643d0ed2778ea1 \ + --hash=sha256:8293f3dea7fc929ef7240796ba231413afa7b68ce38fd21da2995549f5961981 \ + --hash=sha256:8456928655f856c6e1533ff59d5be76578a7157224dbd9ce6872f25055ab9ab7 \ + --hash=sha256:890bcb4abd5a2d3f852196437129eb3667d62630333aacc13dfd470fad3aaa82 \ + --hash=sha256:94a76daa32eb78d61339aff7952ea819b1734b46f73646a07decb40e5b3448e2 \ + --hash=sha256:9f16fbdf4da055efb21c22d81b89f155f02ba420558db21288b3d0035bafd5f4 \ + --hash=sha256:a3d1fae9863299076f05cb8a778c467578262fae09f9dc0ee9b12eb4268ce663 \ + --hash=sha256:a3d507bb6a513ca96ba84443226af944b0f7f47dcc9a399d110cd6146481d24c \ + --hash=sha256:abace499247268e3757271b2f1e244b36b06f8515cf27c4d49468fc9eb16e93d \ + --hash=sha256:ba2a27ff02f48193fc4daeadf8ad2590516fa3d0adeeb34336b96f7fa64c1e3a \ + --hash=sha256:bc84e875994c3b445871ea7181d424588171efec3e185dced958dad9e001950a \ + --hash=sha256:bfd56bb4b37ed4f330b82402f6f435845a5f5648edf1ad497da51a8452d5d62d \ + --hash=sha256:c18ff11e86df2e28854939acde2d003f7984f721eba450b56a200ad90eeb0e6b \ + --hash=sha256:c3bcce8521d785d510b2aad26ae2c966092b7daa8f45dd8f44734a104dc0bc1a \ + --hash=sha256:c4143987a42a2397f2fc3b4d7e3a7d313fbe684f67ff443999e803dd75a76826 \ + --hash=sha256:c69fd885df7d089548a42d5ec05be26050ebcd2283d89b3d30676eb32ff87dee \ + --hash=sha256:ced80795227d70549a411a4ab66e8ce307899fad2220ce5ab2f296e687eacde9 \ + --hash=sha256:d66e421495fdb797610a08f43b05269e0a5ea7f5e652a89bfd5a7d3c1dee3648 \ + --hash=sha256:d861ee9e76ace6cf36a6a89b959ec08e7bc2493ee39d07ffe5acb23ef46d27da \ + --hash=sha256:e9251e3be159d1020c4030bd2e5f84d6a43fe54b6c19c12f51cde9542a2817b2 \ + --hash=sha256:f145bba11b878005c496e93e257c1e88f154d278d2638e6450d17e0f31e558d2 \ + --hash=sha256:fe346b143ff9685e40192a4960938545c699054ba11d4f9029f94751e3f71d87 + # via + # -r requirements_formatting.txt.in + # pyjwt darker==1.7.2 \ --hash=sha256:ec5b7c382d9537611c164f3ecca2e1b8a7923bc5a02bf22f6e7f6c8bcbdf593a \ --hash=sha256:ec9d130ab2a0f7fa49ab68a08fd231a5bec66147ecbbf94c92a1f33d97b5ef6f @@ -274,18 +311,35 @@ pyjwt[crypto]==2.9.0 \ --hash=sha256:3b02fb0f44517787776cf48f2ae25d8e14f300e6d7545a4315cee571a415e850 \ --hash=sha256:7e1e5b56cc735432a7369cbfa0efe50fa113ebecdc04ae6922deba8b84582d0c # via pygithub -pynacl==1.5.0 \ - --hash=sha256:06b8f6fa7f5de8d5d2f7573fe8c863c051225a27b61e6860fd047b1775807858 \ - --hash=sha256:0c84947a22519e013607c9be43706dd42513f9e6ae5d39d3613ca1e142fba44d \ - --hash=sha256:20f42270d27e1b6a29f54032090b972d97f0a1b0948cc52392041ef7831fee93 \ - --hash=sha256:401002a4aaa07c9414132aaed7f6836ff98f59277a234704ff66878c2ee4a0d1 \ - --hash=sha256:52cb72a79269189d4e0dc537556f4740f7f0a9ec41c1322598799b0bdad4ef92 \ - --hash=sha256:61f642bf2378713e2c2e1de73444a3778e5f0a38be6fee0fe532fe30060282ff \ - --hash=sha256:8ac7448f09ab85811607bdd21ec2464495ac8b7c66d146bf545b0f08fb9220ba \ - --hash=sha256:a36d4a9dda1f19ce6e03c9a784a2921a4b726b02e1c736600ca9c22029474394 \ - --hash=sha256:a422368fc821589c228f4c49438a368831cb5bbc0eab5ebe1d7fac9dded6567b \ - --hash=sha256:e46dae94e34b085175f8abb3b0aaa7da40767865ac82c928eeb9e57e1ea8a543 - # via pygithub +pynacl==1.6.2 \ + --hash=sha256:018494d6d696ae03c7e656e5e74cdfd8ea1326962cc401bcf018f1ed8436811c \ + --hash=sha256:04316d1fc625d860b6c162fff704eb8426b1a8bcd3abacea11142cbd99a6b574 \ + --hash=sha256:22de65bb9010a725b0dac248f353bb072969c94fa8d6b1f34b87d7953cf7bbe4 \ + --hash=sha256:26bfcd00dcf2cf160f122186af731ae30ab120c18e8375684ec2670dccd28130 \ + --hash=sha256:2fef529ef3ee487ad8113d287a593fa26f48ee3620d92ecc6f1d09ea38e0709b \ + --hash=sha256:320ef68a41c87547c91a8b58903c9caa641ab01e8512ce291085b5fe2fcb7590 \ + --hash=sha256:3bffb6d0f6becacb6526f8f42adfb5efb26337056ee0831fb9a7044d1a964444 \ + --hash=sha256:44081faff368d6c5553ccf55322ef2819abb40e25afaec7e740f159f74813634 \ + --hash=sha256:46065496ab748469cdd999246d17e301b2c24ae2fdf739132e580a0e94c94a87 \ + --hash=sha256:5811c72b473b2f38f7e2a3dc4f8642e3a3e9b5e7317266e4ced1fba85cae41aa \ + --hash=sha256:622d7b07cc5c02c666795792931b50c91f3ce3c2649762efb1ef0d5684c81594 \ + --hash=sha256:62985f233210dee6548c223301b6c25440852e13d59a8b81490203c3227c5ba0 \ + --hash=sha256:68be3a09455743ff9505491220b64440ced8973fe930f270c8e07ccfa25b1f9e \ + --hash=sha256:834a43af110f743a754448463e8fd61259cd4ab5bbedcf70f9dabad1d28a394c \ + --hash=sha256:8845c0631c0be43abdd865511c41eab235e0be69c81dc66a50911594198679b0 \ + --hash=sha256:8a66d6fb6ae7661c58995f9c6435bda2b1e68b54b598a6a10247bfcdadac996c \ + --hash=sha256:8b097553b380236d51ed11356c953bf8ce36a29a3e596e934ecabe76c985a577 \ + --hash=sha256:a84bf1c20339d06dc0c85d9aea9637a24f718f375d861b2668b2f9f96fa51145 \ + --hash=sha256:a9f9932d8d2811ce1a8ffa79dcbdf3970e7355b5c8eb0c1a881a57e7f7d96e88 \ + --hash=sha256:bc4a36b28dd72fb4845e5d8f9760610588a96d5a51f01d84d8c6ff9849968c14 \ + --hash=sha256:c8a231e36ec2cab018c4ad4358c386e36eede0319a0c41fed24f840b1dac59f6 \ + --hash=sha256:c949ea47e4206af7c8f604b8278093b674f7c79ed0d4719cc836902bf4517465 \ + --hash=sha256:d071c6a9a4c94d79eb665db4ce5cedc537faf74f2355e4d502591d850d3913c0 \ + --hash=sha256:d29bfe37e20e015a7d8b23cfc8bd6aa7909c92a1b8f41ee416bbb3e79ef182b2 \ + --hash=sha256:fe9847ca47d287af41e82be1dd5e23023d3c31a951da134121ab02e42ac218c9 + # via + # -r requirements_formatting.txt.in + # pygithub requests==2.32.3 \ --hash=sha256:55365417734eb18255590a9ff9eb97e9e1da868d4ccd6402399eaf68af20a760 \ --hash=sha256:70761cfe03c773ceb22aa2f671b4757976145175cdfca038c02654d061d6dcc6 @@ -294,10 +348,67 @@ toml==0.10.2 \ --hash=sha256:806143ae5bfb6a3c6e736a764057db0e6a0e05e338b5630894a5f779cabb4f9b \ --hash=sha256:b3bda1d108d5dd99f4a20d24d9c348e91c4db7ab1b749200bded2f839ccbe68f # via darker -urllib3==2.2.2 \ - --hash=sha256:a448b2f64d686155468037e1ace9f2d2199776e17f0a46610480d311f73e3472 \ - --hash=sha256:dd505485549a7a552833da5e6063639d0d177c04f23bc3864e41e5dc5f612168 - # via requests +tomli==2.4.0 \ + --hash=sha256:0408e3de5ec77cc7f81960c362543cbbd91ef883e3138e81b729fc3eea5b9729 \ + --hash=sha256:0dc56fef0e2c1c470aeac5b6ca8cc7b640bb93e92d9803ddaf9ea03e198f5b0b \ + --hash=sha256:0e0fe8a0b8312acf3a88077a0802565cb09ee34107813bba1c7cd591fa6cfc8d \ + --hash=sha256:0f2e3955efea4d1cfbcb87bc321e00dc08d2bcb737fd1d5e398af111d86db5df \ + --hash=sha256:133e93646ec4300d651839d382d63edff11d8978be23da4cc106f5a18b7d0576 \ + --hash=sha256:1b168f2731796b045128c45982d3a4874057626da0e2ef1fdd722848b741361d \ + --hash=sha256:1c8a885b370751837c029ef9bc014f27d80840e48bac415f3412e6593bbc18c1 \ + --hash=sha256:1f776e7d669ebceb01dee46484485f43a4048746235e683bcdffacdf1fb4785a \ + --hash=sha256:1fb2945cbe303b1419e2706e711b7113da57b7db31ee378d08712d678a34e51e \ + --hash=sha256:20cedb4ee43278bc4f2fee6cb50daec836959aadaf948db5172e776dd3d993fc \ + --hash=sha256:20ffd184fb1df76a66e34bd1b36b4a4641bd2b82954befa32fe8163e79f1a702 \ + --hash=sha256:26ab906a1eb794cd4e103691daa23d95c6919cc2fa9160000ac02370cc9dd3f6 \ + --hash=sha256:2add28aacc7425117ff6364fe9e06a183bb0251b03f986df0e78e974047571fd \ + --hash=sha256:2b1e3b80e1d5e52e40e9b924ec43d81570f0e7d09d11081b797bc4692765a3d4 \ + --hash=sha256:31d556d079d72db7c584c0627ff3a24c5d3fb4f730221d3444f3efb1b2514776 \ + --hash=sha256:36b9d05b51e65b254ea6c2585b59d2c4cb91c8a3d91d0ed0f17591a29aaea54a \ + --hash=sha256:39b0b5d1b6dd03684b3fb276407ebed7090bbec989fa55838c98560c01113b66 \ + --hash=sha256:3cf226acb51d8f1c394c1b310e0e0e61fecdd7adcb78d01e294ac297dd2e7f87 \ + --hash=sha256:3d895d56bd3f82ddd6faaff993c275efc2ff38e52322ea264122d72729dca2b2 \ + --hash=sha256:413540dce94673591859c4c6f794dfeaa845e98bf35d72ed59636f869ef9f86f \ + --hash=sha256:43e685b9b2341681907759cf3a04e14d7104b3580f808cfde1dfdb60ada85475 \ + --hash=sha256:4cbcb367d44a1f0c2be408758b43e1ffb5308abe0ea222897d6bfc8e8281ef2f \ + --hash=sha256:551e321c6ba03b55676970b47cb1b73f14a0a4dce6a3e1a9458fd6d921d72e95 \ + --hash=sha256:5572e41282d5268eb09a697c89a7bee84fae66511f87533a6f88bd2f7b652da9 \ + --hash=sha256:5aa48d7c2356055feef06a43611fc401a07337d5b006be13a30f6c58f869e3c3 \ + --hash=sha256:5b5807f3999fb66776dbce568cc9a828544244a8eb84b84b9bafc080c99597b9 \ + --hash=sha256:5e3f639a7a8f10069d0e15408c0b96a2a828cfdec6fca05296ebcdcc28ca7c76 \ + --hash=sha256:685306e2cc7da35be4ee914fd34ab801a6acacb061b6a7abca922aaf9ad368da \ + --hash=sha256:75c2f8bbddf170e8effc98f5e9084a8751f8174ea6ccf4fca5398436e0320bc8 \ + --hash=sha256:7b438885858efd5be02a9a133caf5812b8776ee0c969fea02c45e8e3f296ba51 \ + --hash=sha256:7d49c66a7d5e56ac959cb6fc583aff0651094ec071ba9ad43df785abc2320d86 \ + --hash=sha256:7d6d9a4aee98fac3eab4952ad1d73aee87359452d1c086b5ceb43ed02ddb16b8 \ + --hash=sha256:84d081fbc252d1b6a982e1870660e7330fb8f90f676f6e78b052ad4e64714bf0 \ + --hash=sha256:8768715ffc41f0008abe25d808c20c3d990f42b6e2e58305d5da280ae7d1fa3b \ + --hash=sha256:920b1de295e72887bafa3ad9f7a792f811847d57ea6b1215154030cf131f16b1 \ + --hash=sha256:9a08144fa4cba33db5255f9b74f0b89888622109bd2776148f2597447f92a94e \ + --hash=sha256:a26d7ff68dfdb9f87a016ecfd1e1c2bacbe3108f4e0f8bcd2228ef9a766c787d \ + --hash=sha256:aa89c3f6c277dd275d8e243ad24f3b5e701491a860d5121f2cdd399fbb31fc9c \ + --hash=sha256:b5ef256a3fd497d4973c11bf142e9ed78b150d36f5773f1ca6088c230ffc5867 \ + --hash=sha256:b6c78bdf37764092d369722d9946cb65b8767bfa4110f902a1b2542d8d173c8a \ + --hash=sha256:bbb1b10aa643d973366dc2cb1ad94f99c1726a02343d43cbc011edbfac579e7c \ + --hash=sha256:c084ad935abe686bd9c898e62a02a19abfc9760b5a79bc29644463eaf2840cb0 \ + --hash=sha256:c73add4bb52a206fd0c0723432db123c0c75c280cbd67174dd9d2db228ebb1b4 \ + --hash=sha256:cae9c19ed12d4e8f3ebf46d1a75090e4c0dc16271c5bce1c833ac168f08fb614 \ + --hash=sha256:d20b797a5c1ad80c516e41bc1fb0443ddb5006e9aaa7bda2d71978346aeb9132 \ + --hash=sha256:d3d1654e11d724760cdb37a3d7691f0be9db5fbdaef59c9f532aabf87006dbaa \ + --hash=sha256:d878f2a6707cc9d53a1be1414bbb419e629c3d6e67f69230217bb663e76b5087 + # via black +typing-extensions==4.15.0 \ + --hash=sha256:0cea48d173cc12fa28ecabc3b837ea3cf6f38c6d1136f85cbaaf598984861466 \ + --hash=sha256:f0fa19c6845758ab08074a0cfa8b7aecb71c999ca73d62883bc25cc018c4e548 + # via + # black + # cryptography +urllib3==2.6.3 \ + --hash=sha256:1b62b6884944a57dbe321509ab94fd4d3b307075e0c2eae991ac71ee15ad38ed \ + --hash=sha256:bf272323e553dfb2e87d9bfd225ca7b0f467b919d7bbd355436d3fd37cb0acd4 + # via + # -r requirements_formatting.txt.in + # requests wrapt==1.16.0 \ --hash=sha256:0d2691979e93d06a95a26257adb7bfd0c93818e89b1406f5a28f36e0d8c1e1fc \ --hash=sha256:14d7dc606219cdd7405133c713f2c218d4252f2a469003f8c46bb92d5d095d81 \ diff --git a/llvm/utils/git/requirements_formatting.txt.in b/llvm/utils/git/requirements_formatting.txt.in index 4aac571af1cf5..c6f31b9b444f4 100644 --- a/llvm/utils/git/requirements_formatting.txt.in +++ b/llvm/utils/git/requirements_formatting.txt.in @@ -1,3 +1,6 @@ black~=23.0 darker==1.7.2 PyGithub==1.59.1 +urllib3~=2.6.3 +pynacl~=1.6.2 +cryptography~=46.0.5 From 25c7c3bf60f85231f7b8ea6b635496620407e21d Mon Sep 17 00:00:00 2001 From: Cyndy Ishida Date: Wed, 11 Feb 2026 12:13:24 -0800 Subject: [PATCH 012/705] [clang][modules] Remove fno-modules-check-relocated as a driver option (#180995) This preprocessor option is only ever respected as a CC1 option because it's never handled in the driver. One gets warnings like `argument unused during compilation: '-fno-modules-check-relocated'` when not passed directly to CC1. Move the option definition into the Preprocessor CC1 only visibility section. --- clang/include/clang/Options/Options.td | 9 +++++---- clang/test/Driver/unsupported-option.c | 4 ++++ 2 files changed, 9 insertions(+), 4 deletions(-) diff --git a/clang/include/clang/Options/Options.td b/clang/include/clang/Options/Options.td index 40251a65f8f70..653c103983853 100644 --- a/clang/include/clang/Options/Options.td +++ b/clang/include/clang/Options/Options.td @@ -3645,10 +3645,6 @@ defm implicit_modules : BoolFOption<"implicit-modules", NegFlag, PosFlag, BothFlags< [NoXarchOption], [ClangOption, CLOption]>>; -def fno_modules_check_relocated : Joined<["-"], "fno-modules-check-relocated">, - Group, Visibility<[ClangOption, CC1Option]>, - HelpText<"Skip checks for relocated modules when loading PCM files">, - MarshallingInfoNegativeFlag>; def fretain_comments_from_system_headers : Flag<["-"], "fretain-comments-from-system-headers">, Group, Visibility<[ClangOption, CC1Option]>, MarshallingInfoFlag>; @@ -8988,6 +8984,11 @@ def disable_pragma_debug_crash : Flag<["-"], "disable-pragma-debug-crash">, def source_date_epoch : Separate<["-"], "source-date-epoch">, MetaVarName<"