diff --git a/third_party/llvm/generated.patch b/third_party/llvm/generated.patch index f8fa2d147648e..943d396a43efd 100644 --- a/third_party/llvm/generated.patch +++ b/third_party/llvm/generated.patch @@ -1,360 +1,419 @@ Auto generated patch. Do not edit or delete it, even if empty. -diff -ruN --strip-trailing-cr a/libcxx/include/__memory/array_cookie.h b/libcxx/include/__memory/array_cookie.h ---- a/libcxx/include/__memory/array_cookie.h -+++ b/libcxx/include/__memory/array_cookie.h -@@ -24,7 +24,7 @@ - _LIBCPP_BEGIN_NAMESPACE_STD +diff -ruN --strip-trailing-cr a/clang/test/SemaCXX/pragma-pack-packed-2.cpp b/clang/test/SemaCXX/pragma-pack-packed-2.cpp +--- a/clang/test/SemaCXX/pragma-pack-packed-2.cpp ++++ b/clang/test/SemaCXX/pragma-pack-packed-2.cpp +@@ -1,6 +1,6 @@ + // RUN: %clang_cc1 -triple s390x-ibm-zos -fzos-extensions -fsyntax-only -verify %s + // RUN: %clang_cc1 -triple s390x-ibm-zos -fzos-extensions -fxl-pragma-pack -fsyntax-only -verify %s +-// RUN: %clang -target s390x-ibm-zos -S -emit-llvm -Xclang -verify -fno-xl-pragma-pack %s ++// RUN: %clang -target s390x-ibm-zos -S -emit-llvm -Xclang -verify -fno-xl-pragma-pack -o %t.ll %s - // Trait representing whether a type requires an array cookie at the start of its allocation when --// allocated as `new T[n]` and deallocated as `delete array`. -+// allocated as `new T[n]` and deallocated as `delete[] array`. - // - // Under the Itanium C++ ABI [1], we know that an array cookie is available unless `T` is trivially - // destructible and the call to `operator delete[]` is not a sized operator delete. Under ABIs other -diff -ruN --strip-trailing-cr a/libcxx/include/__memory/unique_ptr.h b/libcxx/include/__memory/unique_ptr.h ---- a/libcxx/include/__memory/unique_ptr.h -+++ b/libcxx/include/__memory/unique_ptr.h -@@ -102,6 +102,12 @@ - }; + #pragma pack(show) // expected-warning {{value of #pragma pack(show) == 8}} + #pragma pack(twobyte) +diff -ruN --strip-trailing-cr a/lld/ELF/Arch/AArch64.cpp b/lld/ELF/Arch/AArch64.cpp +--- a/lld/ELF/Arch/AArch64.cpp ++++ b/lld/ELF/Arch/AArch64.cpp +@@ -1208,10 +1208,12 @@ + } + } - template -+struct __is_default_deleter : false_type {}; -+ -+template -+struct __is_default_deleter > : true_type {}; -+ -+template - struct __unique_ptr_deleter_sfinae { - static_assert(!is_reference<_Deleter>::value, "incorrect specialization"); - typedef const _Deleter& __lval_ref_type; -@@ -307,11 +313,16 @@ - // 1. When an array cookie (see [1]) exists at the beginning of the array allocation, we are - // able to reuse that cookie to extract the size of the array and perform bounds checking. - // An array cookie is a size inserted at the beginning of the allocation by the compiler. --// That size is inserted implicitly when doing `new T[n]` in some cases, and its purpose --// is to allow the runtime to destroy the `n` array elements when doing `delete array`. -+// That size is inserted implicitly when doing `new T[n]` in some cases (as of writing this -+// exactly when the array elements are not trivially destructible), and its main purpose is -+// to allow the runtime to destroy the `n` array elements when doing `delete[] array`. - // When we are able to use array cookies, we reuse information already available in the - // current runtime, so bounds checking does not require changing libc++'s ABI. - // -+// However, note that we cannot assume the presence of an array cookie when a custom deleter -+// is used, because the unique_ptr could have been created from an allocation that wasn't -+// obtained via `new T[n]` (since it may not be deleted with `delete[] arr`). -+// - // 2. When the "bounded unique_ptr" ABI configuration (controlled by `_LIBCPP_ABI_BOUNDED_UNIQUE_PTR`) - // is enabled, we store the size of the allocation (when it is known) so we can check it when - // indexing into the `unique_ptr`. That changes the layout of `std::unique_ptr`, which is -@@ -328,7 +339,7 @@ - // try to fall back to using an array cookie when available. - // - // Finally, note that when this ABI configuration is enabled, we have no choice but to always --// make space for a size to be stored in the unique_ptr. Indeed, while we might want to avoid -+// make space for the size to be stored in the unique_ptr. Indeed, while we might want to avoid - // storing the size when an array cookie is available, knowing whether an array cookie is available - // requires the type stored in the unique_ptr to be complete, while unique_ptr can normally - // accommodate incomplete types. -@@ -339,7 +350,9 @@ - __unique_ptr_array_bounds_stateless() = default; - _LIBCPP_HIDE_FROM_ABI _LIBCPP_CONSTEXPR explicit __unique_ptr_array_bounds_stateless(size_t) {} +-void elf::setAArch64TargetInfo(Ctx &ctx) { ++TargetInfo *elf::getAArch64TargetInfo(Ctx &ctx) { + if ((ctx.arg.andFeatures & GNU_PROPERTY_AARCH64_FEATURE_1_BTI) || +- ctx.arg.zPacPlt) +- ctx.target.reset(new AArch64BtiPac(ctx)); +- else +- ctx.target.reset(new AArch64(ctx)); ++ ctx.arg.zPacPlt) { ++ static AArch64BtiPac t(ctx); ++ return &t; ++ } ++ static AArch64 t(ctx); ++ return &t; + } +diff -ruN --strip-trailing-cr a/lld/ELF/Arch/AMDGPU.cpp b/lld/ELF/Arch/AMDGPU.cpp +--- a/lld/ELF/Arch/AMDGPU.cpp ++++ b/lld/ELF/Arch/AMDGPU.cpp +@@ -219,4 +219,7 @@ + } + } + +-void elf::setAMDGPUTargetInfo(Ctx &ctx) { ctx.target.reset(new AMDGPU(ctx)); } ++TargetInfo *elf::getAMDGPUTargetInfo(Ctx &ctx) { ++ static AMDGPU target(ctx); ++ return ⌖ ++} +diff -ruN --strip-trailing-cr a/lld/ELF/Arch/ARM.cpp b/lld/ELF/Arch/ARM.cpp +--- a/lld/ELF/Arch/ARM.cpp ++++ b/lld/ELF/Arch/ARM.cpp +@@ -1533,7 +1533,10 @@ + "': " + toString(std::move(e))); + } + +-void elf::setARMTargetInfo(Ctx &ctx) { ctx.target.reset(new ARM(ctx)); } ++TargetInfo *elf::getARMTargetInfo(Ctx &ctx) { ++ static ARM target(ctx); ++ return ⌖ ++} -- template ::value, int> = 0> -+ template ::value && __has_array_cookie<_Tp>::value, int> = 0> - _LIBCPP_HIDE_FROM_ABI _LIBCPP_CONSTEXPR bool __in_bounds(_Tp* __ptr, size_t __index) const { - // In constant expressions, we can't check the array cookie so we just pretend that the index - // is in-bounds. The compiler catches invalid accesses anyway. -@@ -349,7 +362,9 @@ - return __index < __cookie; + template void elf::writeARMCmseImportLib(Ctx &); + template void elf::writeARMCmseImportLib(Ctx &); +diff -ruN --strip-trailing-cr a/lld/ELF/Arch/AVR.cpp b/lld/ELF/Arch/AVR.cpp +--- a/lld/ELF/Arch/AVR.cpp ++++ b/lld/ELF/Arch/AVR.cpp +@@ -267,7 +267,10 @@ } + } -- template ::value, int> = 0> -+ template ::value || !__has_array_cookie<_Tp>::value, int> = 0> - _LIBCPP_HIDE_FROM_ABI _LIBCPP_CONSTEXPR bool __in_bounds(_Tp*, size_t) const { - return true; // If we don't have an array cookie, we assume the access is in-bounds +-void elf::setAVRTargetInfo(Ctx &ctx) { ctx.target.reset(new AVR(ctx)); } ++TargetInfo *elf::getAVRTargetInfo(Ctx &ctx) { ++ static AVR target(ctx); ++ return ⌖ ++} + + static uint32_t getEFlags(InputFile *file) { + return cast>(file)->getObj().getHeader().e_flags; +diff -ruN --strip-trailing-cr a/lld/ELF/Arch/Hexagon.cpp b/lld/ELF/Arch/Hexagon.cpp +--- a/lld/ELF/Arch/Hexagon.cpp ++++ b/lld/ELF/Arch/Hexagon.cpp +@@ -404,4 +404,7 @@ } -@@ -365,7 +380,9 @@ - _LIBCPP_HIDE_FROM_ABI _LIBCPP_CONSTEXPR explicit __unique_ptr_array_bounds_stored(size_t __size) : __size_(__size) {} + } - // Use the array cookie if there's one -- template ::value, int> = 0> -+ template ::value && __has_array_cookie<_Tp>::value, int> = 0> - _LIBCPP_HIDE_FROM_ABI _LIBCPP_CONSTEXPR bool __in_bounds(_Tp* __ptr, size_t __index) const { - if (__libcpp_is_constant_evaluated()) - return true; -@@ -374,7 +391,9 @@ +-void elf::setHexagonTargetInfo(Ctx &ctx) { ctx.target.reset(new Hexagon(ctx)); } ++TargetInfo *elf::getHexagonTargetInfo(Ctx &ctx) { ++ static Hexagon target(ctx); ++ return ⌖ ++} +diff -ruN --strip-trailing-cr a/lld/ELF/Arch/LoongArch.cpp b/lld/ELF/Arch/LoongArch.cpp +--- a/lld/ELF/Arch/LoongArch.cpp ++++ b/lld/ELF/Arch/LoongArch.cpp +@@ -893,6 +893,7 @@ } + } + +-void elf::setLoongArchTargetInfo(Ctx &ctx) { +- ctx.target.reset(new LoongArch(ctx)); ++TargetInfo *elf::getLoongArchTargetInfo(Ctx &ctx) { ++ static LoongArch target(ctx); ++ return ⌖ + } +diff -ruN --strip-trailing-cr a/lld/ELF/Arch/Mips.cpp b/lld/ELF/Arch/Mips.cpp +--- a/lld/ELF/Arch/Mips.cpp ++++ b/lld/ELF/Arch/Mips.cpp +@@ -779,23 +779,23 @@ + return cast>(file)->getObj().getHeader().e_flags & EF_MIPS_PIC; + } - // Otherwise, fall back on the stored size (if any) -- template ::value, int> = 0> -+ template ::value || !__has_array_cookie<_Tp>::value, int> = 0> - _LIBCPP_HIDE_FROM_ABI _LIBCPP_CONSTEXPR bool __in_bounds(_Tp*, size_t __index) const { - return __index < __size_; +-void elf::setMipsTargetInfo(Ctx &ctx) { ++TargetInfo *elf::getMipsTargetInfo(Ctx &ctx) { + switch (ctx.arg.ekind) { + case ELF32LEKind: { +- ctx.target.reset(new MIPS(ctx)); +- return; ++ static MIPS t(ctx); ++ return &t; } -@@ -562,7 +581,7 @@ + case ELF32BEKind: { +- ctx.target.reset(new MIPS(ctx)); +- return; ++ static MIPS t(ctx); ++ return &t; } - - _LIBCPP_HIDE_FROM_ABI _LIBCPP_CONSTEXPR_SINCE_CXX23 __add_lvalue_reference_t<_Tp> operator[](size_t __i) const { -- _LIBCPP_ASSERT_VALID_ELEMENT_ACCESS(__checker_.__in_bounds(std::__to_address(__ptr_), __i), -+ _LIBCPP_ASSERT_VALID_ELEMENT_ACCESS(__checker_.__in_bounds(std::__to_address(__ptr_), __i), - "unique_ptr::operator[](index): index out of range"); - return __ptr_[__i]; + case ELF64LEKind: { +- ctx.target.reset(new MIPS(ctx)); +- return; ++ static MIPS t(ctx); ++ return &t; + } + case ELF64BEKind: { +- ctx.target.reset(new MIPS(ctx)); +- return; ++ static MIPS t(ctx); ++ return &t; + } + default: + llvm_unreachable("unsupported target"); +diff -ruN --strip-trailing-cr a/lld/ELF/Arch/MSP430.cpp b/lld/ELF/Arch/MSP430.cpp +--- a/lld/ELF/Arch/MSP430.cpp ++++ b/lld/ELF/Arch/MSP430.cpp +@@ -88,4 +88,7 @@ } -diff -ruN --strip-trailing-cr a/libcxx/test/std/utilities/smartptr/unique.ptr/unique.ptr.class/unique.ptr.observers/assert.subscript.pass.cpp b/libcxx/test/std/utilities/smartptr/unique.ptr/unique.ptr.class/unique.ptr.observers/assert.subscript.pass.cpp ---- a/libcxx/test/std/utilities/smartptr/unique.ptr/unique.ptr.class/unique.ptr.observers/assert.subscript.pass.cpp -+++ b/libcxx/test/std/utilities/smartptr/unique.ptr/unique.ptr.class/unique.ptr.observers/assert.subscript.pass.cpp -@@ -48,31 +48,24 @@ + } - template - void test() { -- // For types with an array cookie, we can always detect OOB accesses. -+ // For types with an array cookie, we can always detect OOB accesses. Note that reliance on an array -+ // cookie is limited to the default deleter, since a unique_ptr with a custom deleter may not have -+ // been allocated with `new T[n]`. - { -- // Check with the default deleter - { -- { -- std::unique_ptr ptr(new WithCookie[5]); -- TEST_LIBCPP_ASSERT_FAILURE(ptr[6], "unique_ptr::operator[](index): index out of range"); -- } -- { -- std::unique_ptr ptr = std::make_unique(5); -- TEST_LIBCPP_ASSERT_FAILURE(ptr[6], "unique_ptr::operator[](index): index out of range"); -- } --#if TEST_STD_VER >= 20 -- { -- std::unique_ptr ptr = std::make_unique_for_overwrite(5); -- TEST_LIBCPP_ASSERT_FAILURE(ptr[6] = WithCookie(), "unique_ptr::operator[](index): index out of range"); -- } --#endif -+ std::unique_ptr ptr(new WithCookie[5]); -+ TEST_LIBCPP_ASSERT_FAILURE(ptr[6], "unique_ptr::operator[](index): index out of range"); - } -- -- // Check with a custom deleter - { -- std::unique_ptr ptr(new WithCookie[5]); -+ std::unique_ptr ptr = std::make_unique(5); - TEST_LIBCPP_ASSERT_FAILURE(ptr[6], "unique_ptr::operator[](index): index out of range"); - } -+#if TEST_STD_VER >= 20 -+ { -+ std::unique_ptr ptr = std::make_unique_for_overwrite(5); -+ TEST_LIBCPP_ASSERT_FAILURE(ptr[6] = WithCookie(), "unique_ptr::operator[](index): index out of range"); -+ } -+#endif +-void elf::setMSP430TargetInfo(Ctx &ctx) { ctx.target.reset(new MSP430(ctx)); } ++TargetInfo *elf::getMSP430TargetInfo(Ctx &ctx) { ++ static MSP430 target(ctx); ++ return ⌖ ++} +diff -ruN --strip-trailing-cr a/lld/ELF/Arch/PPC64.cpp b/lld/ELF/Arch/PPC64.cpp +--- a/lld/ELF/Arch/PPC64.cpp ++++ b/lld/ELF/Arch/PPC64.cpp +@@ -1747,4 +1747,7 @@ + return true; + } + +-void elf::setPPC64TargetInfo(Ctx &ctx) { ctx.target.reset(new PPC64(ctx)); } ++TargetInfo *elf::getPPC64TargetInfo(Ctx &ctx) { ++ static PPC64 target(ctx); ++ return ⌖ ++} +diff -ruN --strip-trailing-cr a/lld/ELF/Arch/PPC.cpp b/lld/ELF/Arch/PPC.cpp +--- a/lld/ELF/Arch/PPC.cpp ++++ b/lld/ELF/Arch/PPC.cpp +@@ -523,4 +523,7 @@ } + } - // For types that don't have an array cookie, things are a bit more complicated. We can detect OOB accesses -@@ -97,14 +90,9 @@ - #endif +-void elf::setPPCTargetInfo(Ctx &ctx) { ctx.target.reset(new PPC(ctx)); } ++TargetInfo *elf::getPPCTargetInfo(Ctx &ctx) { ++ static PPC target(ctx); ++ return ⌖ ++} +diff -ruN --strip-trailing-cr a/lld/ELF/Arch/RISCV.cpp b/lld/ELF/Arch/RISCV.cpp +--- a/lld/ELF/Arch/RISCV.cpp ++++ b/lld/ELF/Arch/RISCV.cpp +@@ -1329,4 +1329,7 @@ + mergeAttributesSection(ctx, sections)); + } + +-void elf::setRISCVTargetInfo(Ctx &ctx) { ctx.target.reset(new RISCV(ctx)); } ++TargetInfo *elf::getRISCVTargetInfo(Ctx &ctx) { ++ static RISCV target(ctx); ++ return ⌖ ++} +diff -ruN --strip-trailing-cr a/lld/ELF/Arch/SPARCV9.cpp b/lld/ELF/Arch/SPARCV9.cpp +--- a/lld/ELF/Arch/SPARCV9.cpp ++++ b/lld/ELF/Arch/SPARCV9.cpp +@@ -193,4 +193,7 @@ + relocateNoSym(buf + 4, R_SPARC_WDISP19, -(off + 4 - pltEntrySize)); + } - // Make sure that we carry the bounds information properly through conversions, assignments, etc. -- // These tests are mostly relevant when the ABI setting is enabled (with a stateful bounds-checker), -- // but we still run them for types with an array cookie either way. -+ // These tests are only relevant when the ABI setting is enabled (with a stateful bounds-checker). - #if defined(_LIBCPP_ABI_BOUNDED_UNIQUE_PTR) -- using Types = types::type_list; --#else -- using Types = types::type_list; --#endif -- types::for_each(Types(), [] { -+ types::for_each(types::type_list(), [] { - // Bounds carried through move construction - { - std::unique_ptr ptr = std::make_unique(5); -@@ -135,6 +123,7 @@ - TEST_LIBCPP_ASSERT_FAILURE(other[6], "unique_ptr::operator[](index): index out of range"); - } - }); -+#endif +-void elf::setSPARCV9TargetInfo(Ctx &ctx) { ctx.target.reset(new SPARCV9(ctx)); } ++TargetInfo *elf::getSPARCV9TargetInfo(Ctx &ctx) { ++ static SPARCV9 target(ctx); ++ return ⌖ ++} +diff -ruN --strip-trailing-cr a/lld/ELF/Arch/SystemZ.cpp b/lld/ELF/Arch/SystemZ.cpp +--- a/lld/ELF/Arch/SystemZ.cpp ++++ b/lld/ELF/Arch/SystemZ.cpp +@@ -600,4 +600,7 @@ + } } - template -diff -ruN --strip-trailing-cr a/libcxx/test/std/utilities/smartptr/unique.ptr/unique.ptr.class/unique.ptr.observers/op_subscript.runtime.pass.cpp b/libcxx/test/std/utilities/smartptr/unique.ptr/unique.ptr.class/unique.ptr.observers/op_subscript.runtime.pass.cpp ---- a/libcxx/test/std/utilities/smartptr/unique.ptr/unique.ptr.class/unique.ptr.observers/op_subscript.runtime.pass.cpp -+++ b/libcxx/test/std/utilities/smartptr/unique.ptr/unique.ptr.class/unique.ptr.observers/op_subscript.runtime.pass.cpp -@@ -46,6 +46,11 @@ - template - struct CustomDeleter : std::default_delete {}; +-void elf::setSystemZTargetInfo(Ctx &ctx) { ctx.target.reset(new SystemZ(ctx)); } ++TargetInfo *elf::getSystemZTargetInfo(Ctx &ctx) { ++ static SystemZ t(ctx); ++ return &t; ++} +diff -ruN --strip-trailing-cr a/lld/ELF/Arch/X86_64.cpp b/lld/ELF/Arch/X86_64.cpp +--- a/lld/ELF/Arch/X86_64.cpp ++++ b/lld/ELF/Arch/X86_64.cpp +@@ -1237,17 +1237,21 @@ + write32le(buf + 8, ctx.in.plt->getVA() - pltEntryAddr - 12); + } -+struct NoopDeleter { -+ template -+ TEST_CONSTEXPR_CXX23 void operator()(T*) const {} -+}; -+ - TEST_CONSTEXPR_CXX23 bool test() { - // Basic test - { -@@ -112,12 +117,33 @@ - WithNonTrivialDtor<16>, - WithNonTrivialDtor<256>>; - types::for_each(TrickyCookieTypes(), [] { -- types::for_each(types::type_list, CustomDeleter>(), [] { -- std::unique_ptr p(new T[3]); -+ // Array allocated with `new T[n]`, default deleter -+ { -+ std::unique_ptr> p(new T[3]); -+ assert(p[0] == T()); -+ assert(p[1] == T()); -+ assert(p[2] == T()); -+ } -+ -+ // Array allocated with `new T[n]`, custom deleter -+ { -+ std::unique_ptr> p(new T[3]); -+ assert(p[0] == T()); -+ assert(p[1] == T()); -+ assert(p[2] == T()); -+ } -+ -+ // Array not allocated with `new T[n]`, custom deleter -+ // -+ // This test aims to ensure that the implementation doesn't try to use an array cookie -+ // when there is none. -+ { -+ T array[50] = {}; -+ std::unique_ptr p(&array[0]); - assert(p[0] == T()); - assert(p[1] == T()); - assert(p[2] == T()); -- }); -+ } - }); +-void elf::setX86_64TargetInfo(Ctx &ctx) { ++TargetInfo *elf::getX86_64TargetInfo(Ctx &ctx) { + if (ctx.arg.zRetpolineplt) { +- if (ctx.arg.zNow) +- ctx.target.reset(new RetpolineZNow(ctx)); +- else +- ctx.target.reset(new Retpoline(ctx)); +- return; ++ if (ctx.arg.zNow) { ++ static RetpolineZNow t(ctx); ++ return &t; ++ } ++ static Retpoline t(ctx); ++ return &t; } - #endif // C++20 -diff -ruN --strip-trailing-cr a/lldb/test/API/functionalities/data-formatter/data-formatter-stl/libcxx-simulators/string/main.cpp b/lldb/test/API/functionalities/data-formatter/data-formatter-stl/libcxx-simulators/string/main.cpp ---- a/lldb/test/API/functionalities/data-formatter/data-formatter-stl/libcxx-simulators/string/main.cpp -+++ b/lldb/test/API/functionalities/data-formatter/data-formatter-stl/libcxx-simulators/string/main.cpp -@@ -20,11 +20,7 @@ - // Pre-D128285 layout. - #define PACKED_ANON_STRUCT - #endif --#if REVISION <= 4 --// Pre-2a1ef74 layout. --#define NON_STANDARD_PADDING --#endif --// REVISION == 5: current layout -+// REVISION == 4: current layout - #ifdef PACKED_ANON_STRUCT - #define BEGIN_PACKED_ANON_STRUCT struct __attribute__((packed)) { -@@ -38,7 +34,6 @@ - namespace std { - namespace __lldb { +- if (ctx.arg.andFeatures & GNU_PROPERTY_X86_FEATURE_1_IBT) +- ctx.target.reset(new IntelIBT(ctx)); +- else +- ctx.target.reset(new X86_64(ctx)); ++ if (ctx.arg.andFeatures & GNU_PROPERTY_X86_FEATURE_1_IBT) { ++ static IntelIBT t(ctx); ++ return &t; ++ } ++ ++ static X86_64 t(ctx); ++ return &t; + } +diff -ruN --strip-trailing-cr a/lld/ELF/Arch/X86.cpp b/lld/ELF/Arch/X86.cpp +--- a/lld/ELF/Arch/X86.cpp ++++ b/lld/ELF/Arch/X86.cpp +@@ -706,17 +706,21 @@ + write32le(buf + 22, -off - 26); + } --#ifdef NON_STANDARD_PADDING - #if defined(ALTERNATE_LAYOUT) && defined(SUBCLASS_PADDING) - template struct __padding { - unsigned char __xx[sizeof(_CharT) - 1]; -@@ -46,13 +41,6 @@ +-void elf::setX86TargetInfo(Ctx &ctx) { ++TargetInfo *elf::getX86TargetInfo(Ctx &ctx) { + if (ctx.arg.zRetpolineplt) { +- if (ctx.arg.isPic) +- ctx.target.reset(new RetpolinePic(ctx)); +- else +- ctx.target.reset(new RetpolineNoPic(ctx)); +- return; ++ if (ctx.arg.isPic) { ++ static RetpolinePic t(ctx); ++ return &t; ++ } ++ static RetpolineNoPic t(ctx); ++ return &t; + } - template struct __padding<_CharT, 1> {}; - #endif --#else // !NON_STANDARD_PADDING --template struct __padding { -- char __padding_[_PaddingSize]; --}; -- --template <> struct __padding<0> {}; --#endif +- if (ctx.arg.andFeatures & GNU_PROPERTY_X86_FEATURE_1_IBT) +- ctx.target.reset(new IntelIBT(ctx)); +- else +- ctx.target.reset(new X86(ctx)); ++ if (ctx.arg.andFeatures & GNU_PROPERTY_X86_FEATURE_1_IBT) { ++ static IntelIBT t(ctx); ++ return &t; ++ } ++ ++ static X86 t(ctx); ++ return &t; + } +diff -ruN --strip-trailing-cr a/lld/ELF/Config.h b/lld/ELF/Config.h +--- a/lld/ELF/Config.h ++++ b/lld/ELF/Config.h +@@ -545,7 +545,7 @@ + Config arg; + LinkerDriver driver; + LinkerScript *script; +- std::unique_ptr target; ++ TargetInfo *target; + + // These variables are initialized by Writer and should not be used before + // Writer is initialized. +diff -ruN --strip-trailing-cr a/lld/ELF/Driver.cpp b/lld/ELF/Driver.cpp +--- a/lld/ELF/Driver.cpp ++++ b/lld/ELF/Driver.cpp +@@ -99,7 +99,7 @@ + driver.~LinkerDriver(); + new (&driver) LinkerDriver(*this); + script = nullptr; +- target.reset(); ++ target = nullptr; - template class basic_string { - public: -@@ -89,12 +77,7 @@ - }; - #else // !SUBCLASS_PADDING + bufferStart = nullptr; + mainPart = nullptr; +@@ -3126,7 +3126,7 @@ + // The Target instance handles target-specific stuff, such as applying + // relocations or writing a PLT section. It also contains target-dependent + // values such as a default image base address. +- setTarget(ctx); ++ ctx.target = getTarget(ctx); --#ifdef NON_STANDARD_PADDING - unsigned char __padding[sizeof(value_type) - 1]; --#else -- [[no_unique_address]] __padding __padding_; --#endif -- - #ifdef BITMASKS - unsigned char __size_; - #else // !BITMASKS -@@ -146,26 +129,21 @@ - union { - #ifdef BITMASKS - unsigned char __size_; --#else // !BITMASKS -+#else - struct { - unsigned char __is_long_ : 1; - unsigned char __size_ : 7; - }; --#endif // BITMASKS -+#endif - value_type __lx; - }; --#else // !SHORT_UNION -+#else - BEGIN_PACKED_ANON_STRUCT - unsigned char __is_long_ : 1; - unsigned char __size_ : 7; - END_PACKED_ANON_STRUCT --#ifdef NON_STANDARD_PADDING -- unsigned char __padding[sizeof(value_type) - 1]; --#else // !NON_STANDARD_PADDING -- [[no_unique_address]] __padding __padding_; --#endif // NON_STANDARD_PADDING -- --#endif // SHORT_UNION -+ char __padding_[sizeof(value_type) - 1]; -+#endif - value_type __data_[__min_cap]; - }; + ctx.arg.eflags = ctx.target->calcEFlags(); + // maxPageSize (sometimes called abi page size) is the maximum page size that +diff -ruN --strip-trailing-cr a/lld/ELF/Target.cpp b/lld/ELF/Target.cpp +--- a/lld/ELF/Target.cpp ++++ b/lld/ELF/Target.cpp +@@ -45,39 +45,39 @@ + return std::string(s); + } -diff -ruN --strip-trailing-cr a/lldb/test/API/functionalities/data-formatter/data-formatter-stl/libcxx-simulators/string/TestDataFormatterLibcxxStringSimulator.py b/lldb/test/API/functionalities/data-formatter/data-formatter-stl/libcxx-simulators/string/TestDataFormatterLibcxxStringSimulator.py ---- a/lldb/test/API/functionalities/data-formatter/data-formatter-stl/libcxx-simulators/string/TestDataFormatterLibcxxStringSimulator.py -+++ b/lldb/test/API/functionalities/data-formatter/data-formatter-stl/libcxx-simulators/string/TestDataFormatterLibcxxStringSimulator.py -@@ -27,7 +27,7 @@ +-void elf::setTarget(Ctx &ctx) { ++TargetInfo *elf::getTarget(Ctx &ctx) { + switch (ctx.arg.emachine) { + case EM_386: + case EM_IAMCU: +- return setX86TargetInfo(ctx); ++ return getX86TargetInfo(ctx); + case EM_AARCH64: +- return setAArch64TargetInfo(ctx); ++ return getAArch64TargetInfo(ctx); + case EM_AMDGPU: +- return setAMDGPUTargetInfo(ctx); ++ return getAMDGPUTargetInfo(ctx); + case EM_ARM: +- return setARMTargetInfo(ctx); ++ return getARMTargetInfo(ctx); + case EM_AVR: +- return setAVRTargetInfo(ctx); ++ return getAVRTargetInfo(ctx); + case EM_HEXAGON: +- return setHexagonTargetInfo(ctx); ++ return getHexagonTargetInfo(ctx); + case EM_LOONGARCH: +- return setLoongArchTargetInfo(ctx); ++ return getLoongArchTargetInfo(ctx); + case EM_MIPS: +- return setMipsTargetInfo(ctx); ++ return getMipsTargetInfo(ctx); + case EM_MSP430: +- return setMSP430TargetInfo(ctx); ++ return getMSP430TargetInfo(ctx); + case EM_PPC: +- return setPPCTargetInfo(ctx); ++ return getPPCTargetInfo(ctx); + case EM_PPC64: +- return setPPC64TargetInfo(ctx); ++ return getPPC64TargetInfo(ctx); + case EM_RISCV: +- return setRISCVTargetInfo(ctx); ++ return getRISCVTargetInfo(ctx); + case EM_SPARCV9: +- return setSPARCV9TargetInfo(ctx); ++ return getSPARCV9TargetInfo(ctx); + case EM_S390: +- return setSystemZTargetInfo(ctx); ++ return getSystemZTargetInfo(ctx); + case EM_X86_64: +- return setX86_64TargetInfo(ctx); ++ return getX86_64TargetInfo(ctx); + default: + fatal("unsupported e_machine value: " + Twine(ctx.arg.emachine)); + } +diff -ruN --strip-trailing-cr a/lld/ELF/Target.h b/lld/ELF/Target.h +--- a/lld/ELF/Target.h ++++ b/lld/ELF/Target.h +@@ -179,21 +179,21 @@ + uint64_t defaultImageBase = 0x10000; + }; +-void setAArch64TargetInfo(Ctx &); +-void setAMDGPUTargetInfo(Ctx &); +-void setARMTargetInfo(Ctx &); +-void setAVRTargetInfo(Ctx &); +-void setHexagonTargetInfo(Ctx &); +-void setLoongArchTargetInfo(Ctx &); +-void setMSP430TargetInfo(Ctx &); +-void setMipsTargetInfo(Ctx &); +-void setPPC64TargetInfo(Ctx &); +-void setPPCTargetInfo(Ctx &); +-void setRISCVTargetInfo(Ctx &); +-void setSPARCV9TargetInfo(Ctx &); +-void setSystemZTargetInfo(Ctx &); +-void setX86TargetInfo(Ctx &); +-void setX86_64TargetInfo(Ctx &); ++TargetInfo *getAArch64TargetInfo(Ctx &); ++TargetInfo *getAMDGPUTargetInfo(Ctx &); ++TargetInfo *getARMTargetInfo(Ctx &); ++TargetInfo *getAVRTargetInfo(Ctx &); ++TargetInfo *getHexagonTargetInfo(Ctx &); ++TargetInfo *getLoongArchTargetInfo(Ctx &); ++TargetInfo *getMSP430TargetInfo(Ctx &); ++TargetInfo *getMipsTargetInfo(Ctx &); ++TargetInfo *getPPC64TargetInfo(Ctx &); ++TargetInfo *getPPCTargetInfo(Ctx &); ++TargetInfo *getRISCVTargetInfo(Ctx &); ++TargetInfo *getSPARCV9TargetInfo(Ctx &); ++TargetInfo *getSystemZTargetInfo(Ctx &); ++TargetInfo *getX86TargetInfo(Ctx &); ++TargetInfo *getX86_64TargetInfo(Ctx &); - for v in [None, "ALTERNATE_LAYOUT"]: -- for r in range(6): -+ for r in range(5): - for c in range(3): - name = "test_r%d_c%d" % (r, c) - defines = ["REVISION=%d" % r, "COMPRESSED_PAIR_REV=%d" % c] -diff -ruN --strip-trailing-cr a/utils/bazel/llvm-project-overlay/libc/BUILD.bazel b/utils/bazel/llvm-project-overlay/libc/BUILD.bazel ---- a/utils/bazel/llvm-project-overlay/libc/BUILD.bazel -+++ b/utils/bazel/llvm-project-overlay/libc/BUILD.bazel -@@ -148,6 +148,9 @@ - libc_support_library( - name = "hdr_stdio_macros", - hdrs = ["hdr/stdio_macros.h"], -+ deps = [ -+ ":hdr_stdio_overlay", -+ ], - ) + struct ErrorPlace { + InputSectionBase *isec; +@@ -251,7 +251,7 @@ + void createTaggedSymbols(Ctx &); + void initSymbolAnchors(Ctx &); - libc_support_library( -@@ -215,11 +218,17 @@ - libc_support_library( - name = "types_off_t", - hdrs = ["hdr/types/off_t.h"], -+ deps = [ -+ ":hdr_stdio_overlay", -+ ], - ) +-void setTarget(Ctx &); ++TargetInfo *getTarget(Ctx &); - libc_support_library( - name = "types_FILE", - hdrs = ["hdr/types/FILE.h"], -+ deps = [ -+ ":hdr_stdio_overlay", -+ ], - ) + template bool isMipsPIC(const Defined *sym); - libc_support_library( diff --git a/third_party/llvm/workspace.bzl b/third_party/llvm/workspace.bzl index 50e3160cadca3..e939a2ea30102 100644 --- a/third_party/llvm/workspace.bzl +++ b/third_party/llvm/workspace.bzl @@ -4,8 +4,8 @@ load("//third_party:repo.bzl", "tf_http_archive") def repo(name): """Imports LLVM.""" - LLVM_COMMIT = "82f5acfbec65e1a645d902f746253eeaf0bd2d70" - LLVM_SHA256 = "c621acf6f202cfcbd0c73f381272b70c51c95f0ecb62c17d2d5fe0156a80082a" + LLVM_COMMIT = "2918e779a9545a66c0031b03b3af5bf4d8517cec" + LLVM_SHA256 = "efd6ac33aba00a6eac0047533a99d6d766dc9dae9ec5ba8e74ff297bfd4c9f0e" tf_http_archive( name = name, diff --git a/third_party/shardy/temporary.patch b/third_party/shardy/temporary.patch index 39ab6038ffb2a..23eb28a2ac232 100644 --- a/third_party/shardy/temporary.patch +++ b/third_party/shardy/temporary.patch @@ -1,2670 +1,791 @@ -diff --git a/shardy/dialect/sdy/ir/attrs.td b/shardy/dialect/sdy/ir/attrs.td -index e133c5c..ff8a885 100644 ---- a/shardy/dialect/sdy/ir/attrs.td -+++ b/shardy/dialect/sdy/ir/attrs.td -@@ -95,8 +95,6 @@ def Sdy_Mesh : AttrDef { - ]; - - let extraClassDeclaration = [{ -- bool hasAxis(StringRef axisName) const; -- - int64_t getAxisSize(StringRef axisName) const; - - // Returns the total size of the mesh across all axes, as in the total -@@ -363,43 +361,28 @@ def Sdy_DimensionSharding : AttrDef { - }]; - } - --// Either a `MeshAttr` or a symbol name, referencing a corresponding `MeshOp` --// symbol. --def Sdy_MeshOrRef : AnyAttrOf<[Sdy_Mesh, FlatSymbolRefAttr], -- "mesh attr or flat mesh symbol reference attr"> { -- string cppType = "::mlir::Attribute"; --} -- - def Sdy_TensorSharding : AttrDef { - let mnemonic = "sharding"; - let summary = "Tensor sharding"; - let description = [{ -- A tensor sharding is bound to a specific mesh, and can only reference axis -- names from that mesh. The dimension shardings tell us for each dimension of -- the tensor, along which axes (or sub-axes) it is sharded from major to -- minor. All other axes that don’t shard a dimension are either implicitly or -- explicitly (if they appear in the list of replicated axes) replicated. -- -- The mesh this sharding is bound to can either be specified by a symbol -- name, referencing a corresponding `MeshOp` symbol, or an inlined `MeshAttr`. -+ A tensor sharding is bound to a specific mesh by its name, and can only -+ reference axis names from that mesh. The dimension shardings tell us for -+ each dimension of the tensor, along which axes (or sub-axes) it is sharded -+ from major to minor. All other axes that don’t shard a dimension are either -+ implicitly or explicitly (if they appear in the list of replicated axes) -+ replicated. - }]; - let parameters = (ins -- Sdy_MeshOrRef:$mesh_or_ref, -+ "FlatSymbolRefAttr":$mesh_sym_name, - OptionalArrayRefParameter<"DimensionShardingAttr">:$dim_shardings, - Sdy_AxisRefs:$replicated_axes - ); - let assemblyFormat = [{ -- `<` custom($mesh_or_ref) `,` `[` (`]`):($dim_shardings^ `]`)? `` -+ `<` $mesh_sym_name `,` `[` (`]`):($dim_shardings^ `]`)? `` - (`,` `replicated` `` `=` `` `{` $replicated_axes^ `}`)? `>` - }]; - - let builders = [ -- AttrBuilder<(ins "StringAttr":$mesh_name, -- "ArrayRef":$dim_shardings, -- "ArrayRef":$replicated_axes), [{ -- return $_get($_ctxt, FlatSymbolRefAttr::get(mesh_name), -- dim_shardings, replicated_axes); -- }]>, - AttrBuilder<(ins "StringRef":$mesh_name, - "ArrayRef":$dim_shardings, - "ArrayRef":$replicated_axes), [{ -@@ -435,29 +418,10 @@ def Sdy_TensorSharding : AttrDef { - }); - } - -- // Returns the mesh `FlatSymbolRefAttr` this sharding references, assuming -- // it doesn't have an inlined `MeshAttr`. -- FlatSymbolRefAttr getMeshSymName() const { -- return mlir::cast(getMeshOrRef()); -- } -- -- // Returns the mesh name this sharding references, assuming it doesn't have -- // an inlined `MeshAttr`. - StringRef getMeshName() const { - return getMeshSymName().getValue(); - } - -- // If this sharding has an inlined `MeshAttr`, returns it, otherwise looks -- // up the mesh symbol with the referenced name in `symbolTable`, and returns -- // its `MeshAttr` if it exists in the table, or nullptr otherwise. -- MeshAttr getMesh(const SymbolTable& symbolTable) const; -- -- // If this sharding has an inlined `MeshAttr`, returns it, otherwise looks -- // up the mesh symbol with the referenced name in the symbol table of the -- // enclosing module of `op`, and returns its `MeshAttr` if it exists in the -- // table, or nullptr otherwise. -- MeshAttr getMesh(Operation* op) const; -- - // Returns true if all dimension shardings are empty and there are no - // replicated axes. - bool emptyAxes() const; -diff --git a/shardy/dialect/sdy/ir/canonicalization.cc b/shardy/dialect/sdy/ir/canonicalization.cc -index 3473b4b..82c0990 100644 ---- a/shardy/dialect/sdy/ir/canonicalization.cc -+++ b/shardy/dialect/sdy/ir/canonicalization.cc -@@ -15,10 +15,10 @@ limitations under the License. - - #include - #include -+#include - - #include "llvm/ADT/STLExtras.h" - #include "mlir/IR/BuiltinAttributes.h" --#include "mlir/IR/Diagnostics.h" - #include "mlir/IR/MLIRContext.h" - #include "mlir/IR/OperationSupport.h" - #include "mlir/IR/PatternMatch.h" -@@ -80,26 +80,21 @@ class RedundantManualComputationPattern - private: - LogicalResult matchAndRewrite(ManualComputationOp manualComputationOp, - PatternRewriter& rewriter) const override { -- ArrayRef inShardings = -- manualComputationOp.getInShardings().getShardings(); -- ArrayRef outShardings = -- manualComputationOp.getOutShardings().getShardings(); -+ std::optional meshName = -+ getCommonMeshName(manualComputationOp.getInShardings().getShardings(), -+ manualComputationOp.getOutShardings().getShardings()); - - int64_t manualAxesProduct = 1; -- if (!inShardings.empty() || !outShardings.empty()) { -- MeshAttr mesh = -- getCommonMesh(inShardings, outShardings, manualComputationOp); -- assert(mesh && "expected inputs and outputs to have a common mesh"); -+ if (meshName.has_value()) { -+ MeshAttr mesh = getMeshAttr(manualComputationOp, *meshName); -+ assert(mesh && "unknown mesh"); - for (StringAttr manualAxis : manualComputationOp.getManualAxes()) { - manualAxesProduct *= mesh.getAxisSize(manualAxis); - } - } - - if (manualAxesProduct != 1) { -- return rewriter.notifyMatchFailure( -- manualComputationOp, [](Diagnostic& diag) { -- diag << "product of manual axis sizes is not 1"; -- }); -+ return failure(); - } - - mlir::InlinerInterface inliner(manualComputationOp.getContext()); -diff --git a/shardy/dialect/sdy/ir/dialect.cc b/shardy/dialect/sdy/ir/dialect.cc -index ef96407..07a0975 100644 ---- a/shardy/dialect/sdy/ir/dialect.cc -+++ b/shardy/dialect/sdy/ir/dialect.cc -@@ -27,7 +27,6 @@ limitations under the License. - #include "llvm/ADT/STLExtras.h" - #include "llvm/ADT/SmallVector.h" - #include "llvm/Support/ErrorHandling.h" --#include "mlir/IR/Attributes.h" - #include "mlir/IR/BuiltinAttributes.h" - #include "mlir/IR/BuiltinTypeInterfaces.h" - #include "mlir/IR/BuiltinTypes.h" -@@ -35,7 +34,6 @@ limitations under the License. - #include "mlir/IR/OpDefinition.h" - #include "mlir/IR/OperationSupport.h" - #include "mlir/IR/Region.h" --#include "mlir/IR/SymbolTable.h" - #include "mlir/IR/TypeRange.h" - #include "mlir/IR/Types.h" - #include "mlir/IR/Value.h" -@@ -159,12 +157,6 @@ void ShardableDataFlowOpInterface::setEdgeOwnerSharding( - // MeshAttr - //===----------------------------------------------------------------------===// - --bool MeshAttr::hasAxis(StringRef axisName) const { -- return llvm::any_of(getAxes(), [axisName](MeshAxisAttr axis) { -- return axis.getName() == axisName; -- }); --} -- - int64_t MeshAttr::getAxisSize(StringRef axisName) const { - for (MeshAxisAttr meshAxis : getAxes()) { - if (meshAxis.getName() == axisName) { -@@ -410,40 +402,17 @@ namespace { - - // Creates fully open or closed tensor sharding attr. - TensorShardingAttr getTensorShardingAttr(MLIRContext* context, int64_t rank, -- Attribute meshOrRef, bool isClosed) { -+ StringRef meshName, bool isClosed) { - return TensorShardingAttr::get( -- context, meshOrRef, -+ context, meshName, - /*dimShardings=*/ - SmallVector( - rank, DimensionShardingAttr::get(context, {}, isClosed)), - /*replicatedAxes=*/{}); - } - --// Creates fully open or closed tensor sharding attr. --TensorShardingAttr getTensorShardingAttr(MLIRContext* context, int64_t rank, -- StringRef meshName, bool isClosed) { -- return getTensorShardingAttr( -- context, rank, FlatSymbolRefAttr::get(context, meshName), isClosed); --} -- - } // namespace - --MeshAttr TensorShardingAttr::getMesh(const SymbolTable& symbolTable) const { -- // TODO(tomnatan): remove mlir:: once Attribute::dyn_cast is removed. -- if (auto mesh = mlir::dyn_cast(getMeshOrRef())) { -- return mesh; -- } -- return getMeshAttr(symbolTable, getMeshSymName()); --} -- --MeshAttr TensorShardingAttr::getMesh(Operation* op) const { -- // TODO(tomnatan): remove mlir:: once Attribute::dyn_cast is removed. -- if (auto mesh = mlir::dyn_cast(getMeshOrRef())) { -- return mesh; -- } -- return getMeshAttr(op, getMeshSymName()); --} -- - bool TensorShardingAttr::emptyAxes() const { - return getReplicatedAxes().empty() && - llvm::all_of(getDimShardings(), -@@ -492,7 +461,7 @@ TensorShardingAttr TensorShardingAttr::closeShardingDims( - dimShardings[dim] = DimensionShardingAttr::get( - getContext(), dimShardings[dim].getAxes(), /*isClosed=*/true); - } -- return TensorShardingAttr::get(getContext(), getMeshOrRef(), dimShardings, -+ return TensorShardingAttr::get(getContext(), getMeshName(), dimShardings, - getReplicatedAxes()); - } - -@@ -504,7 +473,7 @@ TensorShardingAttr TensorShardingAttr::openShardingDims( - dimShardings[dim] = DimensionShardingAttr::get( - getContext(), dimShardings[dim].getAxes(), /*isClosed=*/false); - } -- return TensorShardingAttr::get(getContext(), getMeshOrRef(), dimShardings, -+ return TensorShardingAttr::get(getContext(), getMeshName(), dimShardings, - getReplicatedAxes()); - } - -@@ -512,7 +481,7 @@ TensorShardingAttr TensorShardingAttr::replaceDimSharding( - int64_t dim, DimensionShardingAttr sharding) const { - SmallVector shardings(getDimShardings()); - shardings[dim] = sharding; -- return TensorShardingAttr::get(getContext(), getMeshOrRef(), shardings, -+ return TensorShardingAttr::get(getContext(), getMeshName(), shardings, - getReplicatedAxes()); - } - -@@ -534,8 +503,8 @@ TensorShardingAttr TensorShardingAttr::getReplicated(StringRef axisName, - AxisRefAttr::getMeshComparator(mesh)), - newAxisRef); - -- return TensorShardingAttr::get(getContext(), getMeshOrRef(), -- getDimShardings(), newReplicatedAxes); -+ return TensorShardingAttr::get(getContext(), getMeshName(), getDimShardings(), -+ newReplicatedAxes); - } - - TensorShardingAttr TensorShardingAttr::getFullyClosed(MLIRContext* context, -@@ -552,8 +521,8 @@ TensorShardingAttr TensorShardingAttr::getFullyOpen(MLIRContext* context, - - TensorShardingAttr TensorShardingAttr::getFullyOpenLike( - TensorShardingAttr sharding) { -- return getTensorShardingAttr(sharding.getContext(), sharding.getRank(), -- sharding.getMeshOrRef(), /*isClosed=*/false); -+ return TensorShardingAttr::getFullyOpen( -+ sharding.getContext(), sharding.getRank(), sharding.getMeshName()); - } - - //===----------------------------------------------------------------------===// -@@ -720,7 +689,7 @@ TensorShardingAttr addFreeAxesToManualComputationSharding( - newDimSharding.getIsClosed()); - } - return TensorShardingAttr::get( -- returnedSharding.getContext(), returnedSharding.getMeshOrRef(), -+ returnedSharding.getContext(), returnedSharding.getMeshName(), - resultDimShardings, returnedSharding.getReplicatedAxes()); - } - -diff --git a/shardy/dialect/sdy/ir/parsers.cc b/shardy/dialect/sdy/ir/parsers.cc -index 7fcbdb0..c48bb2a 100644 ---- a/shardy/dialect/sdy/ir/parsers.cc -+++ b/shardy/dialect/sdy/ir/parsers.cc -@@ -104,27 +104,6 @@ Attribute DimensionShardingAttr::parse(AsmParser& parser, Type type) { - priority == -1 ? std::nullopt : std::make_optional(priority)); - } - --ParseResult parseMeshOrRef(AsmParser& parser, Attribute& meshOrRef) { -- if (!parser.parseOptionalKeyword(MeshAttr::getMnemonic())) { -- auto mesh = FieldParser::parse(parser); -- if (failed(mesh)) { -- return parser.emitError(parser.getCurrentLocation(), -- "failed to parse MeshAttr"); -- } -- meshOrRef = *mesh; -- return success(); -- } -- -- auto symbolRef = FieldParser::parse(parser); -- if (failed(symbolRef)) { -- return parser.emitError(parser.getCurrentLocation(), -- "expecting MeshAttr or FlatSymbolRefAttr, got: ") -- << meshOrRef; -- } -- meshOrRef = *symbolRef; -- return success(); --} -- - namespace { - - // Removes and returns the index from the symbol in factorsStr. For example: -diff --git a/shardy/dialect/sdy/ir/parsers.h b/shardy/dialect/sdy/ir/parsers.h -index e4bc9a0..b7fedbd 100644 ---- a/shardy/dialect/sdy/ir/parsers.h -+++ b/shardy/dialect/sdy/ir/parsers.h -@@ -18,7 +18,6 @@ limitations under the License. - - #include - --#include "mlir/IR/Attributes.h" - #include "mlir/IR/OpImplementation.h" - #include "mlir/IR/Region.h" - #include "mlir/Support/LLVM.h" -@@ -28,8 +27,6 @@ limitations under the License. - namespace mlir { - namespace sdy { - --ParseResult parseMeshOrRef(AsmParser& parser, Attribute& meshOrRef); -- - // Parses the factor sizes of an OpShardingRule. The keys in the list are the - // indices in the factor mapping, with i=0, j=1, k=2,... z=17. For any index - // greater than 17 it should be parsed as z_X where `X>0` and the index -diff --git a/shardy/dialect/sdy/ir/printers.cc b/shardy/dialect/sdy/ir/printers.cc -index f4c8031..8be1771 100644 ---- a/shardy/dialect/sdy/ir/printers.cc -+++ b/shardy/dialect/sdy/ir/printers.cc -@@ -18,7 +18,6 @@ limitations under the License. - #include - - #include "llvm/ADT/STLExtras.h" --#include "mlir/IR/Attributes.h" - #include "mlir/IR/BuiltinAttributes.h" - #include "mlir/IR/OpImplementation.h" - #include "mlir/IR/Region.h" -@@ -53,15 +52,6 @@ void DimMappingAttr::print(AsmPrinter& printer) const { - } - } - --void printMeshOrRef(AsmPrinter& printer, Attribute meshOrRef) { -- if (auto mesh = dyn_cast(meshOrRef)) { -- printer << MeshAttr::getMnemonic(); -- printer.printStrippedAttrOrType(mesh); -- } else { -- printer << meshOrRef; -- } --} -- - void printFactorSizes(AsmPrinter& printer, ArrayRef factorSizes) { - if (factorSizes.empty()) { - return; -diff --git a/shardy/dialect/sdy/ir/printers.h b/shardy/dialect/sdy/ir/printers.h -index 956a1fa..caa7946 100644 ---- a/shardy/dialect/sdy/ir/printers.h -+++ b/shardy/dialect/sdy/ir/printers.h -@@ -19,7 +19,6 @@ limitations under the License. - #include - #include - --#include "mlir/IR/Attributes.h" - #include "mlir/IR/OpImplementation.h" - #include "mlir/IR/Operation.h" - #include "mlir/IR/Region.h" -@@ -29,8 +28,6 @@ limitations under the License. - namespace mlir { - namespace sdy { - --void printMeshOrRef(AsmPrinter& printer, Attribute meshOrRef); -- - // Prints the factor sizes of an OpShardingRule. The keys in the list are the - // indices in the factor mapping, with i=0, j=1, k=2,... z=17. For any index - // greater than 17 it's printed as z_X where `X>0` and the index being X-17 -diff --git a/shardy/dialect/sdy/ir/test/manual_computation_canonicalization.mlir b/shardy/dialect/sdy/ir/test/manual_computation_canonicalization.mlir -index d94d232..0d5124f 100644 ---- a/shardy/dialect/sdy/ir/test/manual_computation_canonicalization.mlir -+++ b/shardy/dialect/sdy/ir/test/manual_computation_canonicalization.mlir -@@ -41,17 +41,6 @@ func.func @redundant_manual_axes(%arg0: tensor<8xf32>) -> tensor<8xf32> { - return %0: tensor<8xf32> - } - --// CHECK-LABEL: func @redundant_manual_axes_inlined_mesh --func.func @redundant_manual_axes_inlined_mesh(%arg0: tensor<8xf32>) -> tensor<8xf32> { -- // CHECK-NEXT: %0 = stablehlo.custom_call @foo(%arg0) {has_side_effect = true} : (tensor<8xf32>) -> tensor<8xf32> -- // CHECK-NEXT: return %0 : tensor<8xf32> -- %0 = sdy.manual_computation(%arg0) in_shardings=[, [{"b", "c", "a"}]>] out_shardings=[, [{"b", "c", "a"}]>] -- manual_axes={"b", "c"} (%arg1: tensor<8xf32>) { -- %1 = stablehlo.custom_call @foo(%arg1) {has_side_effect = true} : (tensor<8xf32>) -> tensor<8xf32> -- sdy.return %1 : tensor<8xf32> -- } : (tensor<8xf32>) -> (tensor<8xf32>) -- return %0: tensor<8xf32> --} - - // Top manual_computation is inlinable, middle one isn't, most inner is. - // CHECK-LABEL: func @nested_two_inlinable -diff --git a/shardy/dialect/sdy/ir/test/manual_computation_verification.mlir b/shardy/dialect/sdy/ir/test/manual_computation_verification.mlir -index d43411d..ab7781d 100644 ---- a/shardy/dialect/sdy/ir/test/manual_computation_verification.mlir -+++ b/shardy/dialect/sdy/ir/test/manual_computation_verification.mlir -@@ -1,10 +1,10 @@ - // RUN: sdy_opt %s -split-input-file -verify-diagnostics - - sdy.mesh @meshA = <["a"=2]> --sdy.mesh @meshB = <["a"=4]> -+sdy.mesh @meshB = <["a"=2]> - - func.func @man_comp_different_meshes(%arg0: tensor<16x32xf32>) -> tensor<16x32xf32> { -- // expected-error @+1 {{op result shardings can only be bound to one mesh: #sdy.mesh<["a"=4]> vs #sdy.mesh<["a"=2]>}} -+ // expected-error @+1 {{op result shardings can only refer to one mesh: @meshB vs @meshA}} - %0 = sdy.manual_computation(%arg0) in_shardings=[<@meshA, [{}, {}]>] out_shardings=[<@meshB, [{}, {}]>] manual_axes={} (%arg1: tensor<16x32xf32>) { - %1 = stablehlo.add %arg1, %arg1 : tensor<16x32xf32> - sdy.return %1 : tensor<16x32xf32> -@@ -25,19 +25,6 @@ func.func @unknown_mesh(%arg0: tensor<16x32xf32>) -> tensor<16x32xf32> { - - // ----- - --sdy.mesh @meshA = <["a"=2]> -- --func.func @unknown_manual_axis(%arg0: tensor<16x32xf32>) -> tensor<16x32xf32> { -- // expected-error @+1 {{unknown manual axis: "b"}} -- %0 = sdy.manual_computation(%arg0) in_shardings=[<@meshA, [{"a"}, {}]>] out_shardings=[<@meshA, [{"a"}, {?}]>] manual_axes={"a", "b"} (%arg1: tensor<8x32xf32>) { -- %1 = stablehlo.add %arg1, %arg1 : tensor<8x32xf32> -- sdy.return %1 : tensor<8x32xf32> -- } : (tensor<16x32xf32>) -> tensor<16x32xf32> -- func.return %0: tensor<16x32xf32> --} -- --// ----- -- - func.func @manual_computation_no_inputs_or_outputs_with_manual_axes() { - // expected-error @+1 {{op cannot have manual_axes when there are no input/output shardings.}} - sdy.manual_computation() in_shardings=[] out_shardings=[] manual_axes={"a"} () { -diff --git a/shardy/dialect/sdy/ir/test/tensor_sharding_parse_print.mlir b/shardy/dialect/sdy/ir/test/tensor_sharding_parse_print.mlir -index 0578cf5..5ae2990 100644 ---- a/shardy/dialect/sdy/ir/test/tensor_sharding_parse_print.mlir -+++ b/shardy/dialect/sdy/ir/test/tensor_sharding_parse_print.mlir -@@ -36,22 +36,6 @@ func.func @replicated_axes(%arg0 : tensor<8x8xf32>, %arg1 : tensor<8x8xf32>) -> - return %0 : tensor<8x8xf32> - } - --// CHECK-LABEL: func @inlined_mesh --func.func @inlined_mesh(%arg0 : tensor<8x8xf32>, %arg1 : tensor<8x8xf32>) -> tensor<8x8xf32> { -- // CHECK-NEXT: stablehlo.add -- // CHECK-SAME{LITERAL}: #sdy.sharding_per_value<[, [{"x"}, {}]>]> -- %0 = stablehlo.add %arg0, %arg1 {sdy.sharding = #sdy.sharding_per_value<[, [{"x"}, {}]>]>} : tensor<8x8xf32> -- return %0 : tensor<8x8xf32> --} -- --// CHECK-LABEL: func @inlined_mesh_and_ref --func.func @inlined_mesh_and_ref(%arg0 : tensor<8x8xf32>, %arg1 : tensor<8x8xf32>) -> (tensor<8x8xf32>, tensor<8x8xf32>) { -- // CHECK-NEXT: stablehlo.optimization_barrier -- // CHECK-SAME{LITERAL}: #sdy.sharding_per_value<[, [{"a"}, {}]>, <@foo, [{"a"}, {}]>]> -- %0:2 = stablehlo.optimization_barrier {sdy.sharding = #sdy.sharding_per_value<[, [{"a"}, {}]>, <@foo, [{"a"}, {}]>]>} %arg0, %arg1 : tensor<8x8xf32>, tensor<8x8xf32> -- return %0#0, %0#1 : tensor<8x8xf32>, tensor<8x8xf32> --} -- - // CHECK-LABEL: func @tensor_with_open_dimension_shardings - func.func @tensor_with_open_dimension_shardings(%arg0 : tensor<8x8xf32>, %arg1 : tensor<8x8xf32>) -> tensor<8x8xf32> { - // CHECK-NEXT: stablehlo.add -diff --git a/shardy/dialect/sdy/ir/test/tensor_sharding_verification.mlir b/shardy/dialect/sdy/ir/test/tensor_sharding_verification.mlir -index 098a0ff..a8adf82 100644 ---- a/shardy/dialect/sdy/ir/test/tensor_sharding_verification.mlir -+++ b/shardy/dialect/sdy/ir/test/tensor_sharding_verification.mlir -@@ -115,10 +115,10 @@ func.func @num_shardings_does_not_match_num_results(%arg0: tensor<2x64x13xf32>, - sdy.mesh @mesh1 = <["a"=2]> - sdy.mesh @mesh2 = <["b"=2]> - --// CHECK-LABEL: func @op_shardings_refer_to_different_meshes --func.func @op_shardings_refer_to_different_meshes(%arg0: tensor<2x64x13xf32>, %arg1: tensor<2x64x13xf32>) -> (tensor<2x13xf32>, tensor<2x13xf32>) { -+// CHECK-LABEL: func @op_shardings_refers_to_multiples_meshes -+func.func @op_shardings_refers_to_multiples_meshes(%arg0: tensor<2x64x13xf32>, %arg1: tensor<2x64x13xf32>) -> (tensor<2x13xf32>, tensor<2x13xf32>) { - %0 = stablehlo.constant dense<0.000000e+00> : tensor -- // expected-error @+1 {{op result shardings can only be bound to one mesh: #sdy.mesh<["b"=2]> vs #sdy.mesh<["a"=2]>}} -+ // expected-error @+1 {{op result shardings can only refer to one mesh: @mesh2 vs @mesh1}} - %1:2 = stablehlo.reduce(%arg0 init: %0), (%arg1 init: %0) across dimensions = [1] - {sdy.sharding=#sdy.sharding_per_value<[<@mesh1, [{"a"}, {}]>, <@mesh2, [{"b"}, {}]>]>} : - (tensor<2x64x13xf32>, tensor<2x64x13xf32>, tensor, tensor) -> (tensor<2x13xf32>, tensor<2x13xf32>) -@@ -132,42 +132,6 @@ func.func @op_shardings_refer_to_different_meshes(%arg0: tensor<2x64x13xf32>, %a - - // ----- - --// CHECK-LABEL: func @op_shardings_bound_to_different_inlined_meshes --func.func @op_shardings_bound_to_different_inlined_meshes(%arg0: tensor<2x64x13xf32>, %arg1: tensor<2x64x13xf32>) -> (tensor<2x13xf32>, tensor<2x13xf32>) { -- %0 = stablehlo.constant dense<0.000000e+00> : tensor -- // expected-error @+1 {{op result shardings can only be bound to one mesh: #sdy.mesh<["b"=2]> vs #sdy.mesh<["a"=2]>}} -- %1:2 = stablehlo.reduce(%arg0 init: %0), (%arg1 init: %0) across dimensions = [1] -- {sdy.sharding=#sdy.sharding_per_value<[, [{"a"}, {}]>, , [{"b"}, {}]>]>} : -- (tensor<2x64x13xf32>, tensor<2x64x13xf32>, tensor, tensor) -> (tensor<2x13xf32>, tensor<2x13xf32>) -- reducer(%arg2: tensor, %arg4: tensor) (%arg3: tensor, %arg5: tensor) { -- %2 = stablehlo.add %arg2, %arg4 : tensor -- %3 = stablehlo.add %arg3, %arg5 : tensor -- stablehlo.return %2, %3 : tensor, tensor -- } -- return %1#0, %1#1 : tensor<2x13xf32>, tensor<2x13xf32> --} -- --// ----- -- --sdy.mesh @mesh2 = <["b"=2]> -- --// CHECK-LABEL: func @op_shardings_bound_to_different_inlined_and_referenced_meshes --func.func @op_shardings_bound_to_different_inlined_and_referenced_meshes(%arg0: tensor<2x64x13xf32>, %arg1: tensor<2x64x13xf32>) -> (tensor<2x13xf32>, tensor<2x13xf32>) { -- %0 = stablehlo.constant dense<0.000000e+00> : tensor -- // expected-error @+1 {{op result shardings can only be bound to one mesh: #sdy.mesh<["b"=2]> vs #sdy.mesh<["a"=2]>}} -- %1:2 = stablehlo.reduce(%arg0 init: %0), (%arg1 init: %0) across dimensions = [1] -- {sdy.sharding=#sdy.sharding_per_value<[, [{"a"}, {}]>, <@mesh2, [{"b"}, {}]>]>} : -- (tensor<2x64x13xf32>, tensor<2x64x13xf32>, tensor, tensor) -> (tensor<2x13xf32>, tensor<2x13xf32>) -- reducer(%arg2: tensor, %arg4: tensor) (%arg3: tensor, %arg5: tensor) { -- %2 = stablehlo.add %arg2, %arg4 : tensor -- %3 = stablehlo.add %arg3, %arg5 : tensor -- stablehlo.return %2, %3 : tensor, tensor -- } -- return %1#0, %1#1 : tensor<2x13xf32>, tensor<2x13xf32> --} -- --// ----- -- - sdy.mesh @mesh = <["a"=2]> - - // The purpose of this test is to check the error msg prefix for a func arg. -@@ -227,7 +191,7 @@ func.func @multi_result_op_failure(%arg0: tensor<2x64x13xf32>, %arg1: tensor<2x6 - - sdy.mesh @mesh = <["a"=2]> - --func.func @unknown_axis_mesh_ref(%arg0: tensor<8x8xf32>, %arg1: tensor<8x8xf32>) -> tensor<8x8xf32> { -+func.func @unknown_axis(%arg0: tensor<8x8xf32>, %arg1: tensor<8x8xf32>) -> tensor<8x8xf32> { - // expected-error @+1 {{unknown axis name: "c"}} - %0 = stablehlo.add %arg0, %arg1 {sdy.sharding=#sdy.sharding_per_value<[<@mesh, [{}, {"c"}]>]>} : tensor<8x8xf32> - return %0 : tensor<8x8xf32> -@@ -235,14 +199,6 @@ func.func @unknown_axis_mesh_ref(%arg0: tensor<8x8xf32>, %arg1: tensor<8x8xf32>) - - // ----- - --func.func @unknown_axis_inlined_mesh(%arg0: tensor<8x8xf32>, %arg1: tensor<8x8xf32>) -> tensor<8x8xf32> { -- // expected-error @+1 {{unknown axis name: "c"}} -- %0 = stablehlo.add %arg0, %arg1 {sdy.sharding=#sdy.sharding_per_value<[, [{}, {"c"}]>]>} : tensor<8x8xf32> -- return %0 : tensor<8x8xf32> --} -- --// ----- -- - sdy.mesh @mesh = <["a"=2]> - - func.func @unknown_replicated_axis(%arg0: tensor<8x8xf32>, %arg1: tensor<8x8xf32>) -> tensor<8x8xf32> { -diff --git a/shardy/dialect/sdy/ir/utils.cc b/shardy/dialect/sdy/ir/utils.cc -index 41d118b..4f07365 100644 ---- a/shardy/dialect/sdy/ir/utils.cc -+++ b/shardy/dialect/sdy/ir/utils.cc -@@ -28,7 +28,6 @@ limitations under the License. - #include "mlir/Dialect/Func/IR/FuncOps.h" - #include "mlir/IR/Attributes.h" - #include "mlir/IR/BuiltinAttributes.h" --#include "mlir/IR/BuiltinOps.h" - #include "mlir/IR/BuiltinTypeInterfaces.h" - #include "mlir/IR/Diagnostics.h" - #include "mlir/IR/MLIRContext.h" -@@ -152,42 +151,19 @@ MeshAttr getMeshAttr(Operation* op, SymbolRefAttr meshSymName) { - return nullptr; - } - --MeshAttr getCommonMesh(ArrayRef operandShardings, -- ArrayRef resultsShardings, -- Operation* op) { -- SymbolTable symbolTable(op->getParentOfType()); -- MeshAttr mesh; -- for (TensorShardingAttr sharding : llvm::concat( -- operandShardings, resultsShardings)) { -- if (!sharding) { -- continue; -- } -- MeshAttr otherMesh = sharding.getMesh(symbolTable); -- if (!mesh) { -- mesh = otherMesh; -- } else if (otherMesh != mesh) { -- // Found more than one mesh name. -- return nullptr; -- } -- } -- -- return mesh; --} -- - std::optional getCommonMeshName( - ArrayRef operandShardings, - ArrayRef resultsShardings) { - StringRef meshName; - for (TensorShardingAttr sharding : llvm::concat( - operandShardings, resultsShardings)) { -- if (!sharding) { -- continue; -- } -- if (meshName.empty()) { -- meshName = sharding.getMeshName(); -- } else if (meshName != sharding.getMeshName()) { -- // Found more than one mesh name. -- return std::nullopt; -+ if (sharding) { -+ if (meshName.empty()) { -+ meshName = sharding.getMeshName(); -+ } else if (meshName != sharding.getMeshName()) { -+ // Found more than one mesh name. -+ return std::nullopt; -+ } - } - } - -@@ -338,15 +314,6 @@ void setSharding(Value value, TensorShardingAttr sharding) { - }); - } - --TensorShardingAttr getFuncResultSharding(FuncOp funcOp, int64_t resNum) { -- return funcOp.getResultAttrOfType(resNum, kShardingAttr); --} -- --void setFuncResultSharding(FuncOp funcOp, int64_t resNum, -- TensorShardingAttr sharding) { -- funcOp.setResultAttr(resNum, kShardingAttr, sharding); --} -- - SmallVector getShardings(ValueRange values) { - return llvm::to_vector( - llvm::map_range(values, [](Value value) { return getSharding(value); })); -@@ -364,9 +331,6 @@ TensorShardingPerValueAttr getShardingPerValue(Operation* op) { - } - - void setShardings(Operation* op, ArrayRef shardings) { -- if (shardings.empty()) { -- return; -- } - setShardings(op, - TensorShardingPerValueAttr::get(op->getContext(), shardings)); - } -diff --git a/shardy/dialect/sdy/ir/utils.h b/shardy/dialect/sdy/ir/utils.h -index d1a3fdd..dc44038 100644 ---- a/shardy/dialect/sdy/ir/utils.h -+++ b/shardy/dialect/sdy/ir/utils.h -@@ -107,12 +107,6 @@ MeshAttr getMeshAttr(Operation* op, StringRef meshName); - // table, or nullptr otherwise. - MeshAttr getMeshAttr(Operation* op, SymbolRefAttr meshSymName); - --// Returns the common `MeshAttr` bound by all the `TensorShardingAttr` or --// nullptr if there is none. --MeshAttr getCommonMesh(ArrayRef operandShardings, -- ArrayRef resultsShardings, -- Operation* op); -- - // Returns the common mesh name used by all the `TensorShardingAttr` or - // std::nullopt if there is none. - std::optional getCommonMeshName( -@@ -193,13 +187,6 @@ void replaceShardingAtIndex(Operation* op, unsigned index, - // TODO(tomnatan): consider moving this to a dedicated sdy/utils dir. - void setSharding(Value value, TensorShardingAttr sharding); - --// Return the sharding of the `resNum` result of the given `funcOp`. --TensorShardingAttr getFuncResultSharding(func::FuncOp funcOp, int64_t resNum); -- --// Sets the sharding of the `resNum` result of the given `funcOp` to `sharding`. --void setFuncResultSharding(func::FuncOp funcOp, int64_t resNum, -- TensorShardingAttr sharding); -- - // Returns the sharding of each value in `values`. Returns an empty array if the - // op has no sharding attributes. - SmallVector getShardings(ValueRange values); -diff --git a/shardy/dialect/sdy/ir/verifiers.cc b/shardy/dialect/sdy/ir/verifiers.cc -index ab6007b..1715f0e 100644 ---- a/shardy/dialect/sdy/ir/verifiers.cc -+++ b/shardy/dialect/sdy/ir/verifiers.cc -@@ -33,7 +33,6 @@ limitations under the License. - #include "mlir/Dialect/Func/IR/FuncOps.h" - #include "mlir/IR/Attributes.h" - #include "mlir/IR/BuiltinAttributes.h" --#include "mlir/IR/BuiltinOps.h" - #include "mlir/IR/BuiltinTypeInterfaces.h" - #include "mlir/IR/BuiltinTypes.h" - #include "mlir/IR/Diagnostics.h" -@@ -76,18 +75,6 @@ EmitErrorFn getEmitValueInRangeErrorFn(EmitErrorFn emitError, int64_t numValues, - }; - } - --template --FailureOr getMeshOrFail(TensorShardingAttr sharding, -- OpOrSymbolTable opOrSymbolTable, -- EmitErrorFn emitError) { -- MeshAttr mesh = sharding.getMesh(opOrSymbolTable); -- if (!mesh) { -- // We can assume the sharding has a mesh symbol name. -- return emitError("unknown mesh: ") << sharding.getMeshSymName(); -- } -- return mesh; --} -- - // Verifies the following for `axisRefs`: - // - // - All axis names are present in `axisNameToSize`. -@@ -346,17 +333,18 @@ LogicalResult verifyTensorShardingAttr(TensorShardingAttr shardingAttr, - return success(); - } - --// Same as the overload above, but in addition verifies that if `shardingAttr` --// references a mesh by name, it's present in the module's symbol table. -+// Same as the overload above, but in addition verifies that the mesh name in -+// `shardingAttr` is present in the module's symbol table. - LogicalResult verifyTensorShardingAttr(TensorShardingAttr shardingAttr, - Type type, Operation* op, - EmitErrorFn emitError) { -- FailureOr mesh = getMeshOrFail(shardingAttr, op, emitError); -- if (failed(mesh)) { -- return failure(); -+ SymbolRefAttr meshName = shardingAttr.getMeshSymName(); -+ MeshAttr mesh = getMeshAttr(op, meshName); -+ if (!mesh) { -+ return emitError("unknown mesh: ") << meshName; - } - -- return verifyTensorShardingAttr(shardingAttr, type, *mesh, emitError, -+ return verifyTensorShardingAttr(shardingAttr, type, mesh, emitError, - /*checkDivisibility=*/false, - getParentManualComputationOps(op)); - } -@@ -365,15 +353,14 @@ LogicalResult verifyTensorShardingAttr(TensorShardingAttr shardingAttr, - // - // - The number of tensor shardings is equal to the number of tensors (size of - // `types`). --// - All shardings must have the same mesh, and it must be equal to `commonMesh` --// if present. -+// - All shardings reference the same mesh name, and that name is present in the -+// module's symbol table. - // - All shardings are valid (see `verifyTensorShardingAttr`). - // TODO(bartchr): relax this to allow different meshes when the op is a dataflow - // op - LogicalResult verifyTensorShardingPerValueAttr( - TensorShardingPerValueAttr shardingPerValueAttr, TypeRange types, -- Operation* op, EmitErrorFn emitError, const SymbolTable& symbolTable, -- MeshAttr commonMesh = MeshAttr()) { -+ Operation* op, EmitErrorFn emitError, SymbolRefAttr meshName) { - ArrayRef shardingsPerValue = - shardingPerValueAttr.getShardings(); - if (shardingsPerValue.size() != types.size()) { -@@ -382,26 +369,29 @@ LogicalResult verifyTensorShardingPerValueAttr( - << " values"; - } - -+ MeshAttr mesh; - for (auto [index, entry] : - llvm::enumerate(llvm::zip(shardingsPerValue, types))) { - EmitErrorFn valueEmitError = - getEmitValueInRangeErrorFn(emitError, types.size(), index); - auto [shardingAttr, resultType] = entry; -- FailureOr otherMesh = -- getMeshOrFail(shardingAttr, symbolTable, valueEmitError); -- if (failed(otherMesh)) { -- return failure(); -- } -- -- if (!commonMesh) { -- commonMesh = *otherMesh; -- } else if (*otherMesh != commonMesh) { -- return emitError("shardings can only be bound to one mesh: ") -- << *otherMesh << " vs " << commonMesh; -+ SymbolRefAttr otherMeshName = shardingAttr.getMeshSymName(); -+ if (!meshName) { -+ meshName = otherMeshName; -+ } else if (meshName != otherMeshName) { -+ return emitError("shardings can only refer to one mesh: ") -+ << otherMeshName << " vs " << meshName; -+ } -+ -+ if (!mesh) { -+ mesh = getMeshAttr(op, meshName); -+ if (!mesh) { -+ return valueEmitError("unknown mesh: ") << meshName; -+ } - } - - if (failed(verifyTensorShardingAttr( -- shardingAttr, resultType, commonMesh, valueEmitError, -+ shardingAttr, resultType, mesh, valueEmitError, - /*checkDivisibility=*/false, getParentManualComputationOps(op)))) { - return failure(); - } -@@ -410,15 +400,6 @@ LogicalResult verifyTensorShardingPerValueAttr( - return success(); - } - --// Same as the overload above, but creates a `SymbolTable`. --LogicalResult verifyTensorShardingPerValueAttr( -- TensorShardingPerValueAttr shardingPerValueAttr, TypeRange types, -- Operation* op, EmitErrorFn emitError, MeshAttr commonMesh = MeshAttr()) { -- return verifyTensorShardingPerValueAttr( -- shardingPerValueAttr, types, op, emitError, -- SymbolTable(op->getParentOfType()), commonMesh); --} -- - // Verifies an attribute of either a function argument or result. - LogicalResult verifyFuncAttribute(FuncOp funcOp, NamedAttribute attr, Type type, - int64_t index, StringRef valueKindStr) { -@@ -709,7 +690,7 @@ ArrayRef::iterator findManualAxisAfterFreeAxis( - // For each set of op values (operands/results) and corresponding sharding for - // each value, verifies: - // 1. the sharding list itself wrt the mesh and `globalTypes`, --// 2. each sharding is bound to the same `mesh`, -+// 2. each sharding is bound to the same 'meshName', - // 3. the in/out shardings are valid w.r.t the corresponding global - // type, - // 4. the number of global and local tensor inputs/outputs of the op region -@@ -725,8 +706,7 @@ ArrayRef::iterator findManualAxisAfterFreeAxis( - template - LogicalResult verifyManualComputationValue( - ManualComputationOp op, GlobalRangeT globalTypes, LocalRangeT localTypes, -- TensorShardingPerValueAttr shardingPerValueAttr, -- const SymbolTable& symbolTable, MeshAttr mesh, -+ TensorShardingPerValueAttr shardingPerValueAttr, SymbolRefAttr meshName, - const llvm::SmallDenseSet& manualAxesSet, - StringRef valueKindStr) { - if (globalTypes.empty() && localTypes.empty() && -@@ -735,7 +715,7 @@ LogicalResult verifyManualComputationValue( - return success(); - } - // 1. Verify the sharding list itself wrt the mesh and `globalTypes`. -- // 2. Verify each sharding is bound to the same 'mesh'. -+ // 2. Verify each sharding is bound to the same 'meshName'. - // 3. Verify the in/out shardings are valid w.r.t the corresponding global - // type. - if (failed(verifyTensorShardingPerValueAttr( -@@ -743,7 +723,7 @@ LogicalResult verifyManualComputationValue( - [op, valueKindStr](StringRef msg) { - return op->emitOpError(valueKindStr) << " " << msg; - }, -- symbolTable, mesh))) { -+ meshName))) { - return failure(); - } - -@@ -757,6 +737,7 @@ LogicalResult verifyManualComputationValue( - << valueKindStr << "s"; - } - -+ MeshAttr mesh = getMeshAttr(op, meshName); - for (auto [valueIndex, valueEntry] : llvm::enumerate(llvm::zip_equal( - globalTypes, localTypes, shardingPerValueAttr.getShardings()))) { - auto [globalType, localType, sharding] = valueEntry; -@@ -827,7 +808,7 @@ mlir::LogicalResult TensorShardingAttr::verifyForType( - - LogicalResult ManualComputationOp::verify() { - ManualAxisToOwner alreadyManualAxes = -- getParentManualComputationOps(getOperation()); -+ getParentManualComputationOps(this->getOperation()); - for (StringRef axisName : getManualAxes()) { - if (alreadyManualAxes.contains(axisName)) { - return emitBoundAxisInManualComputationError( -@@ -836,35 +817,29 @@ LogicalResult ManualComputationOp::verify() { - } - } - -- // Pick any mesh. `verifyManualComputationValue` will verify all in/out -+ // Pick any mesh. verifyManualComputationValue will verify all in/out - // shardings refer to the same mesh. -- MeshAttr mesh; -+ SymbolRefAttr meshName; - if (!getInShardings().empty()) { -- mesh = getInShardings().getShardings().front().getMesh(getOperation()); -+ meshName = getInShardings().getShardings().front().getMeshSymName(); - } else if (!getOutShardings().empty()) { -- mesh = getOutShardings().getShardings().front().getMesh(getOperation()); -- } else if (!getManualAxes().empty()) { -- return emitOpError( -- "cannot have manual_axes when there are no input/output shardings."); -+ meshName = getOutShardings().getShardings().front().getMeshSymName(); - } -- // If a call to `getMesh` returned empty, a failure will be emitted in -- // `verifyManualComputationValue`. - -- for (StringAttr axisName : getManualAxes()) { -- if (mesh && !mesh.hasAxis(axisName)) { -- return emitOpError("unknown manual axis: ") << axisName; -+ if (!getManualAxes().empty()) { -+ if (!meshName) { -+ return emitOpError( -+ "cannot have manual_axes when there are no input/output shardings."); - } - } -- -- SymbolTable symbolTable(getOperation()->getParentOfType()); - llvm::SmallDenseSet manualAxesSet( - getManualAxes().begin(), getManualAxes().end()); - if (failed(verifyManualComputationValue( - *this, getOperandTypes(), getBody().getArgumentTypes(), -- getInShardings(), symbolTable, mesh, manualAxesSet, "operand")) || -+ getInShardings(), meshName, manualAxesSet, "operand")) || - failed(verifyManualComputationValue( - *this, getResultTypes(), getBodyTerminatorOpOperandTypes(*this), -- getOutShardings(), symbolTable, mesh, manualAxesSet, "result"))) { -+ getOutShardings(), meshName, manualAxesSet, "result"))) { - return failure(); - } - -@@ -923,20 +898,20 @@ LogicalResult NamedComputationOp::verify() { - return success(); - } - -- // TODO(pxy): remove this once the `ShardableDataFlowOpInterface` is verified. - // Verify the in/out shardings. - if (inShardings && - failed(verifyTensorShardingPerValueAttr( -- *inShardings, getOperandTypes(), *this, [this](StringRef msg) { -- return emitOpError("in_shardings ") << msg; -- }))) { -+ *inShardings, getOperandTypes(), *this, -+ [this](StringRef msg) { return emitOpError("in_shardings ") << msg; }, -+ /*meshName=*/nullptr))) { - return failure(); - } -- if (outShardings && -- failed(verifyTensorShardingPerValueAttr( -- *outShardings, getResultTypes(), *this, [this](StringRef msg) { -- return emitOpError("out_shardings ") << msg; -- }))) { -+ if (outShardings && failed(verifyTensorShardingPerValueAttr( -+ *outShardings, getResultTypes(), *this, -+ [this](StringRef msg) { -+ return emitOpError("out_shardings ") << msg; -+ }, -+ /*meshName=*/nullptr))) { - return failure(); - } - -@@ -982,9 +957,12 @@ LogicalResult SdyDialect::verifyOperationAttribute(Operation* op, - << "TensorShardingPerValueAttr"; - } - -+ // TODO(tomnatan): verify all relevant tensors are bound to the same mesh -+ - return verifyTensorShardingPerValueAttr( - shardingPerValue, op->getResultTypes(), op, -- [op](StringRef msg) { return op->emitOpError("result ") << msg; }); -+ [op](StringRef msg) { return op->emitOpError("result ") << msg; }, -+ SymbolRefAttr()); - } - - if (attr.getName() == kShardingRuleAttr) { -diff --git a/shardy/dialect/sdy/transforms/common/BUILD b/shardy/dialect/sdy/transforms/common/BUILD -index f907072..a465e48 100644 ---- a/shardy/dialect/sdy/transforms/common/BUILD -+++ b/shardy/dialect/sdy/transforms/common/BUILD -@@ -34,16 +34,3 @@ cc_test( - "@stablehlo//:stablehlo_ops", - ], - ) -- --cc_library( -- name = "sharding_walker", -- srcs = ["sharding_walker.cc"], -- hdrs = ["sharding_walker.h"], -- deps = [ -- "//shardy/dialect/sdy/ir:dialect", -- "@llvm-project//llvm:Support", -- "@llvm-project//mlir:FuncDialect", -- "@llvm-project//mlir:IR", -- "@llvm-project//mlir:Support", -- ], --) -diff --git a/shardy/dialect/sdy/transforms/common/sharding_walker.cc b/shardy/dialect/sdy/transforms/common/sharding_walker.cc -deleted file mode 100644 -index ad0a968..0000000 ---- a/shardy/dialect/sdy/transforms/common/sharding_walker.cc -+++ /dev/null -@@ -1,172 +0,0 @@ --/* Copyright 2024 The Shardy Authors. -- --Licensed under the Apache License, Version 2.0 (the "License"); --you may not use this file except in compliance with the License. --You may obtain a copy of the License at -- -- http://www.apache.org/licenses/LICENSE-2.0 -- --Unless required by applicable law or agreed to in writing, software --distributed under the License is distributed on an "AS IS" BASIS, --WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. --See the License for the specific language governing permissions and --limitations under the License. --==============================================================================*/ -- --#include "shardy/dialect/sdy/transforms/common/sharding_walker.h" -- --#include --#include --#include -- --#include "llvm/ADT/ArrayRef.h" --#include "llvm/ADT/STLExtras.h" --#include "mlir/Dialect/Func/IR/FuncOps.h" --#include "mlir/IR/Operation.h" --#include "mlir/IR/Value.h" --#include "mlir/Support/LLVM.h" --#include "shardy/dialect/sdy/ir/dialect.h" --#include "shardy/dialect/sdy/ir/utils.h" -- --namespace mlir { --namespace sdy { -- --namespace { -- --using func::FuncOp; -- --// Applies `callback` on `sharding` if present, and calls `setShardingFn` on the --// result if `transformShardings` is true. --void processSharding(TensorShardingAttr sharding, bool transformShardings, -- TransformShardingFn callback, -- std::function setShardingFn) { -- if (!sharding) { -- return; -- } -- TensorShardingAttr newSharding = callback(sharding); -- if (transformShardings && newSharding != sharding) { -- setShardingFn(newSharding); -- } --} -- --// Applies `callback` on each sharding in `shardings` if present, and calls --// `setShardingFn` on the results if `transformShardings` is true. --void processShardings( -- ArrayRef shardings, bool transformShardings, -- TransformShardingFn callback, -- std::function)> setShardingsFn) { -- if (shardings.empty() || !transformShardings) { -- llvm::for_each(shardings, callback); -- return; -- } -- -- SmallVector newShardings; -- llvm::transform(shardings, std::back_inserter(newShardings), callback); -- setShardingsFn(newShardings); --} -- --// Same as above but for `TensorShardingPerValueAttr`. --void processShardings( -- TensorShardingPerValueAttr shardings, bool transformShardings, -- TransformShardingFn callback, -- std::function setShardingsFn) { -- return processShardings(shardings.getShardings(), transformShardings, -- callback, -- [&](ArrayRef newShardings) { -- setShardingsFn(TensorShardingPerValueAttr::get( -- shardings.getContext(), newShardings)); -- }); --} -- --void walkShardings(Operation* rootOp, TransformShardingFn callback, -- ConsumeOpFn consumeOpFn, bool transformShardings) { -- rootOp->walk([&](Operation* op) { -- consumeOpFn(op); -- TypeSwitch(op) -- .Case([&](FuncOp funcOp) { -- for (BlockArgument arg : funcOp.getArguments()) { -- processSharding(getSharding(arg), transformShardings, callback, -- [&](TensorShardingAttr newSharding) { -- setSharding(arg, newSharding); -- }); -- } -- for (int64_t resNum = 0; resNum < funcOp.getNumResults(); ++resNum) { -- processSharding( -- getFuncResultSharding(funcOp, resNum), transformShardings, -- callback, [&](TensorShardingAttr newSharding) { -- setFuncResultSharding(funcOp, resNum, newSharding); -- }); -- } -- }) -- .Case( -- [&](ShardableDataFlowOpInterface shardableDataFlowOp) { -- processShardings( -- shardableDataFlowOp.getBlockArgumentEdgeOwnerShardings(), -- transformShardings, callback, -- [&](ArrayRef newShardings) { -- shardableDataFlowOp.setBlockArgumentEdgeOwnerShardings( -- newShardings); -- }); -- processShardings( -- shardableDataFlowOp.getOpResultEdgeOwnerShardings(), -- transformShardings, callback, -- [&](ArrayRef newShardings) { -- shardableDataFlowOp.setOpResultEdgeOwnerShardings( -- newShardings); -- }); -- }) -- .Case( -- [&](ManualComputationOp manualComputationOp) { -- processShardings( -- manualComputationOp.getInShardings(), transformShardings, -- callback, [&](TensorShardingPerValueAttr newShardings) { -- manualComputationOp.setInShardingsAttr(newShardings); -- }); -- processShardings( -- manualComputationOp.getOutShardings(), transformShardings, -- callback, [&](TensorShardingPerValueAttr newShardings) { -- manualComputationOp.setOutShardingsAttr(newShardings); -- }); -- }) -- .Default([&](Operation* op) { -- if (op->getNumResults() == 1) { -- // For ops with a single result, we use `get/setSharding` instead of -- // `get/setShardings`, since the latter only handle ops with an -- // unregistered sharding attribute, to also handle SDY ops like -- // `ShardingConstraintOp`. -- Value result = op->getResult(0); -- processSharding(getSharding(result), transformShardings, callback, -- [&](TensorShardingAttr newSharding) { -- setSharding(result, newSharding); -- }); -- } else { -- processShardings(getShardings(op), transformShardings, callback, -- [&](ArrayRef newShardings) { -- setShardings(op, newShardings); -- }); -- } -- }); -- }); --} -- --} // namespace -- --void walkShardings(Operation* rootOp, ConsumeShardingFn consumeFn, -- ConsumeOpFn consumeOpFn) { -- walkShardings( -- rootOp, -- [consumeFn](TensorShardingAttr sharding) { -- consumeFn(sharding); -- return sharding; -- }, -- consumeOpFn, -- /*transformShardings=*/false); --} -- --void transformShardings(Operation* rootOp, TransformShardingFn transformFn, -- ConsumeOpFn consumeOpFn) { -- walkShardings(rootOp, transformFn, consumeOpFn, /*transformShardings=*/true); --} -- --} // namespace sdy --} // namespace mlir diff --git a/shardy/dialect/sdy/transforms/common/sharding_walker.h b/shardy/dialect/sdy/transforms/common/sharding_walker.h -deleted file mode 100644 -index ab29e34..0000000 +index 59eda54..1d0e48c 100644 --- a/shardy/dialect/sdy/transforms/common/sharding_walker.h -+++ /dev/null -@@ -1,53 +0,0 @@ --/* Copyright 2024 The Shardy Authors. -- --Licensed under the Apache License, Version 2.0 (the "License"); --you may not use this file except in compliance with the License. --You may obtain a copy of the License at -- -- http://www.apache.org/licenses/LICENSE-2.0 -- --Unless required by applicable law or agreed to in writing, software --distributed under the License is distributed on an "AS IS" BASIS, --WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. --See the License for the specific language governing permissions and --limitations under the License. --==============================================================================*/ -- --#ifndef SHARDY_DIALECT_SDY_TRANSFORMS_COMMON_SHARDING_WALKER_H_ --#define SHARDY_DIALECT_SDY_TRANSFORMS_COMMON_SHARDING_WALKER_H_ -- --#include -- --#include "mlir/IR/Operation.h" --#include "shardy/dialect/sdy/ir/dialect.h" -- --namespace mlir { --namespace sdy { -- --using ConsumeOpFn = std::function; --using ConsumeShardingFn = std::function; --using TransformShardingFn = -- std::function; -- --// Walks the given `rootOp` and applies `callback` on any `TensorShardingAttr` --// encountered. --// --// In addition, applies `consumeOpFn` on every encountered op, before consuming --// its shardings. --void walkShardings( -- Operation* rootOp, ConsumeShardingFn consumeFn, -- ConsumeOpFn consumeOpFn = [](Operation*) {}); -- --// Walks the given `rootOp` and replaces any `TensorShardingAttr` encountered --// with the result of applying `callback` on it. --// --// In addition, applies `consumeOpFn` on every encountered op, before --// transforming its shardings. --void transformShardings( -- Operation* rootOp, TransformShardingFn transformFn, -- ConsumeOpFn consumeOpFn = [](Operation*) {}); -- --} // namespace sdy --} // namespace mlir -- --#endif // SHARDY_DIALECT_SDY_TRANSFORMS_COMMON_SHARDING_WALKER_H_ -diff --git a/shardy/dialect/sdy/transforms/export/update_non_divisible_input_output_shardings.cc b/shardy/dialect/sdy/transforms/export/update_non_divisible_input_output_shardings.cc -index 8290a9c..22c4269 100644 ---- a/shardy/dialect/sdy/transforms/export/update_non_divisible_input_output_shardings.cc -+++ b/shardy/dialect/sdy/transforms/export/update_non_divisible_input_output_shardings.cc -@@ -60,7 +60,8 @@ namespace { - // See update_non_divisible_input_output_shardings.mlir for more examples. - TensorShardingAttr getEvenlySharded(TensorShardingAttr sharding, - ShapedType type, func::FuncOp funcOp) { -- MeshAttr mesh = sharding.getMesh(funcOp); -+ StringRef meshName = sharding.getMeshName(); -+ MeshAttr mesh = getMeshAttr(funcOp, meshName); - assert(mesh && "unknown mesh"); - MLIRContext* ctx = funcOp->getContext(); - llvm::SmallVector newDimShardings; -@@ -118,7 +119,7 @@ TensorShardingAttr getEvenlySharded(TensorShardingAttr sharding, - // NOTE: no need to account for replicated axes, since we end with a sharding - // that covers less-than-or-equal amount of axes than we started with. So no - // way the final sharding can use an axis/sub-axis from the replicated axes. -- return TensorShardingAttr::get(ctx, sharding.getMeshOrRef(), newDimShardings, -+ return TensorShardingAttr::get(ctx, meshName, newDimShardings, - sharding.getReplicatedAxes()); - } - -@@ -155,9 +156,12 @@ struct UpdateNonDivisibleInputOutputShardingsPass - // Update results. - updateValueShardings( - funcOp.getResultTypes(), -- [&](int64_t index) { return getFuncResultSharding(funcOp, index); }, -+ [&](int64_t index) { -+ return funcOp.getResultAttrOfType(index, -+ kShardingAttr); -+ }, - [&](int64_t index, TensorShardingAttr sharding) { -- setFuncResultSharding(funcOp, index, sharding); -+ funcOp.setResultAttr(index, kShardingAttr, sharding); - }, - funcOp); - } -diff --git a/shardy/dialect/sdy/transforms/import/BUILD b/shardy/dialect/sdy/transforms/import/BUILD -index 7585d01..eb618f9 100644 ---- a/shardy/dialect/sdy/transforms/import/BUILD -+++ b/shardy/dialect/sdy/transforms/import/BUILD -@@ -40,7 +40,6 @@ cc_library( - "constant_splitter.cc", - "import_maximal_sharding.cc", - "import_pipeline.cc", -- "lift_inlined_meshes.cc", - "manual_axes_cleanup.cc", - "sharding_group_import.cc", - ], -@@ -52,7 +51,6 @@ cc_library( - "//shardy/common:file_utils", - "//shardy/dialect/sdy/ir:dialect", - "//shardy/dialect/sdy/transforms/common:op_properties", -- "//shardy/dialect/sdy/transforms/common:sharding_walker", - "@llvm-project//llvm:Support", - "@llvm-project//mlir:FuncDialect", - "@llvm-project//mlir:IR", -diff --git a/shardy/dialect/sdy/transforms/import/import_pipeline.cc b/shardy/dialect/sdy/transforms/import/import_pipeline.cc -index 76dc577..7c89a85 100644 ---- a/shardy/dialect/sdy/transforms/import/import_pipeline.cc -+++ b/shardy/dialect/sdy/transforms/import/import_pipeline.cc -@@ -33,7 +33,6 @@ void addImportPipeline(OpPassManager& pm, StringRef dumpDirectory) { - // functions. - pm.addPass(createInlinerPass()); - pm.addPass(createSymbolDCEPass()); -- pm.addPass(createLiftInlinedMeshesPass()); - pm.addNestedPass(createConstantSplitterPass()); - pm.addNestedPass(createAddDataFlowEdgesPass()); - pm.addNestedPass(createApplyShardingConstraintsPass()); -diff --git a/shardy/dialect/sdy/transforms/import/lift_inlined_meshes.cc b/shardy/dialect/sdy/transforms/import/lift_inlined_meshes.cc -deleted file mode 100644 -index 14bcbb6..0000000 ---- a/shardy/dialect/sdy/transforms/import/lift_inlined_meshes.cc -+++ /dev/null -@@ -1,112 +0,0 @@ --/* Copyright 2024 The Shardy Authors. -- --Licensed under the Apache License, Version 2.0 (the "License"); --you may not use this file except in compliance with the License. --You may obtain a copy of the License at -- -- http://www.apache.org/licenses/LICENSE-2.0 -- --Unless required by applicable law or agreed to in writing, software --distributed under the License is distributed on an "AS IS" BASIS, --WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. --See the License for the specific language governing permissions and --limitations under the License. --==============================================================================*/ -- --#include --#include // IWYU pragma: keep -- --#include "llvm/ADT/DenseMap.h" --#include "llvm/ADT/STLExtras.h" --#include "mlir/IR/Attributes.h" --#include "mlir/IR/Builders.h" --#include "mlir/IR/BuiltinAttributes.h" --#include "mlir/IR/BuiltinOps.h" --#include "mlir/IR/SymbolTable.h" --#include "mlir/Pass/Pass.h" // IWYU pragma: keep --#include "mlir/Support/LLVM.h" --#include "shardy/dialect/sdy/ir/dialect.h" --#include "shardy/dialect/sdy/transforms/common/sharding_walker.h" -- --namespace mlir { --namespace sdy { -- --#define GEN_PASS_DEF_LIFTINLINEDMESHESPASS --#include "shardy/dialect/sdy/transforms/import/passes.h.inc" -- --namespace { -- --TensorShardingAttr replaceMesh(TensorShardingAttr sharding, -- StringAttr meshName) { -- return TensorShardingAttr::get(sharding.getContext(), meshName, -- sharding.getDimShardings(), -- sharding.getReplicatedAxes()); --} -- --struct LiftInlinedMeshesPass -- : public impl::LiftInlinedMeshesPassBase { -- using LiftInlinedMeshesPassBase::LiftInlinedMeshesPassBase; -- -- void runOnOperation() final { -- ModuleOp moduleOp = getOperation(); -- SymbolTable symbolTable(moduleOp); -- -- // A map from both: -- // 1. `FlatSymbolRefAttr` referencing a `MeshOp` that was deduped. -- // 2. `MeshAttr` in an existing `MeshOp` or inlined in a -- // `TensorShardingAttr`. -- // To the name of an existing or new `MeshOp` that should now be referenced -- // instead. -- llvm::SmallDenseMap meshOrRefToNewName; -- -- MeshOp lastMeshOp; -- for (auto meshOp : llvm::make_early_inc_range(moduleOp.getOps())) { -- if (auto insertedIt = meshOrRefToNewName.try_emplace( -- meshOp.getMesh(), meshOp.getSymNameAttr()); -- !insertedIt.second) { -- // This is a duplicate mesh, we map its name (as a FlatSymbolRefAttr) to -- // the name of the identical mesh that was already inserted, so we can -- // replace the former with the latter in any sharding in the module that -- // referenced it. We can also erase the mesh because we know it won't be -- // used after this pass. -- // NOTE: assigning to a map entry a value that is read from the same map -- // can lead to use-after-free (if rehash is triggered). -- StringAttr newMeshName = insertedIt.first->second; -- meshOrRefToNewName[FlatSymbolRefAttr::get(meshOp.getSymNameAttr())] = -- newMeshName; -- symbolTable.erase(meshOp); -- } else { -- lastMeshOp = meshOp; -- } -- } -- -- OpBuilder builder(moduleOp); -- if (lastMeshOp) { -- builder.setInsertionPointAfter(lastMeshOp); -- } else { -- builder.setInsertionPointToStart(moduleOp.getBody()); -- } -- -- transformShardings(moduleOp, [&](TensorShardingAttr sharding) { -- // Taking `newMeshName` by reference so we can update it if `sharding` -- // has an inlined mesh that isn't already in the map. -- StringAttr& newMeshName = meshOrRefToNewName[sharding.getMeshOrRef()]; -- if (newMeshName) { -- return replaceMesh(sharding, newMeshName); -- } -- if (auto mesh = dyn_cast(sharding.getMeshOrRef())) { -- // Inlined mesh with a new `MeshAttr`. -- // TODO(tomnatan): give more elaborate names like `maximal_mesh_i`. -- newMeshName = symbolTable.insert( -- builder.create(moduleOp.getLoc(), "mesh", mesh)); -- return replaceMesh(sharding, newMeshName); -- } -- return sharding; -- }); -- } --}; -- --} // namespace -- --} // namespace sdy --} // namespace mlir -diff --git a/shardy/dialect/sdy/transforms/import/manual_axes_cleanup.cc b/shardy/dialect/sdy/transforms/import/manual_axes_cleanup.cc -index 80cecec..2b720e3 100644 ---- a/shardy/dialect/sdy/transforms/import/manual_axes_cleanup.cc -+++ b/shardy/dialect/sdy/transforms/import/manual_axes_cleanup.cc -@@ -24,7 +24,6 @@ limitations under the License. - #include "llvm/ADT/SmallVector.h" - #include "mlir/IR/Builders.h" - #include "mlir/IR/BuiltinAttributes.h" --#include "mlir/IR/Operation.h" - #include "mlir/Pass/Pass.h" // IWYU pragma: keep - #include "mlir/Rewrite/FrozenRewritePatternSet.h" - #include "mlir/Support/LLVM.h" -@@ -41,13 +40,13 @@ namespace { - - // Returns whether the new shardings needed to be updated to account for unused - // manual axes, and the new shardings. --std::optional> --addUnusedManualAxesToReplicatedAxes( -+std::optional> addUnusedManualAxesToReplicatedAxes( - ArrayRef shardings, ArrayRef manualAxes, -- std::function meshComparator) { -+ std::function meshComparator, -+ StringAttr meshName) { - SmallVector newShardings; - bool modified = false; -- for (TensorShardingAttr sharding : shardings) { -+ for (const auto& sharding : shardings) { - llvm::SmallSet unusedManualAxes; - unusedManualAxes.insert(manualAxes.begin(), manualAxes.end()); - sharding.forEachAxisRef([&unusedManualAxes](AxisRefAttr axis) { -@@ -68,7 +67,7 @@ addUnusedManualAxesToReplicatedAxes( - }); - llvm::sort(newReplicatedAxes, meshComparator); - newShardings.push_back( -- TensorShardingAttr::get(sharding.getContext(), sharding.getMeshOrRef(), -+ TensorShardingAttr::get(sharding.getContext(), meshName, - sharding.getDimShardings(), newReplicatedAxes)); - modified = true; - } -@@ -77,29 +76,44 @@ addUnusedManualAxesToReplicatedAxes( - - // Adds any unused manual axes to the replicated_axes list for each in/out - // sharding --void addUnusedManualAxesToReplicatedAxes(ManualComputationOp op, -- MeshAttr mesh) { -+void addUnusedManualAxesToReplicatedAxes(ManualComputationOp op) { - OpBuilder builder(op); -+ std::optional meshName = getCommonMeshName( -+ op.getInShardings().getShardings(), op.getOutShardings().getShardings()); -+ assert(meshName.has_value() && -+ "Expected no input/output ManualComputationOps to be canonicalized"); -+ MeshAttr mesh = getMeshAttr(op, *meshName); -+ assert(mesh && "unknown mesh"); - - ArrayRef manualAxes = op.getManualAxes(); - auto meshComparator = AxisRefAttr::getMeshComparator(mesh); -+ SmallVector newShardings; -+ StringAttr meshNameAttr = builder.getStringAttr(*meshName); - - if (std::optional> newShardings = - addUnusedManualAxesToReplicatedAxes( -- op.getInShardings().getShardings(), manualAxes, meshComparator)) { -+ op.getInShardings().getShardings(), manualAxes, meshComparator, -+ meshNameAttr)) { - op.setInShardings(*newShardings); - } - - if (std::optional> newShardings = - addUnusedManualAxesToReplicatedAxes( -- op.getOutShardings().getShardings(), manualAxes, -- meshComparator)) { -+ op.getOutShardings().getShardings(), manualAxes, meshComparator, -+ meshNameAttr)) { - op.setOutShardings(*newShardings); - } - } - - // Sorts the manual axes in mesh axis order. --void sortManualAxes(ManualComputationOp op, MeshAttr mesh) { -+void sortManualAxes(ManualComputationOp op) { -+ std::optional meshName = getCommonMeshName( -+ op.getInShardings().getShardings(), op.getOutShardings().getShardings()); -+ assert(meshName.has_value() && -+ "Expected no input/output ManualComputationOps to be canonicalized"); -+ MeshAttr mesh = getMeshAttr(op, meshName.value()); -+ assert(mesh && "unknown mesh"); -+ - ArrayRef manualAxes = op.getManualAxes(); - auto meshComparator = mesh.getAxisNameComparator(); - -@@ -117,18 +131,8 @@ struct ManualAxesCleanupPass - - void runOnOperation() final { - getOperation()->walk([](ManualComputationOp op) { -- ArrayRef inShardings = -- op.getInShardings().getShardings(); -- ArrayRef outShardings = -- op.getOutShardings().getShardings(); -- if (inShardings.empty() && outShardings.empty()) { -- // Nothing to do. -- return; -- } -- MeshAttr mesh = getCommonMesh(inShardings, outShardings, op); -- assert(mesh && "expected inputs and outputs to have a common mesh"); -- sortManualAxes(op, mesh); -- addUnusedManualAxesToReplicatedAxes(op, mesh); -+ sortManualAxes(op); -+ addUnusedManualAxesToReplicatedAxes(op); - }); - } - -diff --git a/shardy/dialect/sdy/transforms/import/passes.td b/shardy/dialect/sdy/transforms/import/passes.td -index 547c3a7..1df48dd 100644 ---- a/shardy/dialect/sdy/transforms/import/passes.td -+++ b/shardy/dialect/sdy/transforms/import/passes.td -@@ -15,20 +15,6 @@ limitations under the License. - - include "mlir/Pass/PassBase.td" - --def LiftInlinedMeshesPass : Pass<"sdy-lift-inlined-meshes", "ModuleOp"> { -- let summary = "Lifts inlined `MeshAttr`s in shardings as symbol `MeshOp`s."; -- let description = [{ -- Replaces any inlined `MeshAttr` in a `TensorShardingAttr` with a mesh symbol -- name, referencing either an existing or new `MeshOp` in the module, such -- that no two `MeshOp`s with an identical `MeshAttr` (existing `MeshOp`s are -- deduped as well). -- -- Each new mesh will have the first available name in -- [`mesh`, `mesh_0`, `mesh_1`, ...]. -- }]; -- let dependentDialects = ["mlir::sdy::SdyDialect"]; --} -- - def AddDataFlowEdgesPass : Pass<"sdy-add-data-flow-edges", "func::FuncOp"> { - let summary = "Inserts `DataFlowEdgeOp` for every data-flow edge."; - let description = [{ -diff --git a/shardy/dialect/sdy/transforms/import/test/lift_inlined_meshes.mlir b/shardy/dialect/sdy/transforms/import/test/lift_inlined_meshes.mlir -deleted file mode 100644 -index 1c252e3..0000000 ---- a/shardy/dialect/sdy/transforms/import/test/lift_inlined_meshes.mlir -+++ /dev/null -@@ -1,174 +0,0 @@ --// RUN: sdy_opt -split-input-file %s -sdy-lift-inlined-meshes | FileCheck %s -- --// CHECK: sdy.mesh @mesh1 = <["x"=2, "y"=2]> --sdy.mesh @mesh1 = <["x"=2, "y"=2]> -- --// CHECK: sdy.mesh @mesh2 = <["x"=2, "y"=4]> --sdy.mesh @mesh2 = <["x"=2, "y"=4]> -- --// CHECK-LABEL: func @no_inlined_meshes_or_duplicates( --// CHECK-SAME: %arg0: tensor<8x8xf32> {sdy.sharding = #sdy.sharding<@mesh1, [{"x"}, {}]>}, --// CHECK-SAME: %arg1: tensor<8x8xf32>) --func.func @no_inlined_meshes_or_duplicates(%arg0: tensor<8x8xf32> {sdy.sharding = #sdy.sharding<@mesh1, [{"x"}, {}]>}, %arg1: tensor<8x8xf32>) -> tensor<8x8xf32> { -- // CHECK-NEXT: stablehlo.add %arg0, %arg1 {sdy.sharding = #sdy.sharding_per_value<[<@mesh2, [{"x"}, {}]>]> -- %0 = stablehlo.add %arg0, %arg1 {sdy.sharding = #sdy.sharding_per_value<[<@mesh2, [{"x"}, {}]>]>} : tensor<8x8xf32> -- return %0 : tensor<8x8xf32> --} -- --// ----- -- --// CHECK: sdy.mesh @mesh1 = <["x"=2, "y"=2]> --sdy.mesh @mesh1 = <["x"=2, "y"=2]> -- --// CHECK: sdy.mesh @mesh2 = <["x"=2, "y"=4]> --sdy.mesh @mesh2 = <["x"=2, "y"=4]> -- --// CHECK-NOT: sdy.mesh @mesh3 --sdy.mesh @mesh3 = <["x"=2, "y"=2]> -- --// CHECK-LABEL: func @no_inlined_meshes_with_duplicates( --// CHECK-SAME: %arg0: tensor<8x8xf32> {sdy.sharding = #sdy.sharding<@mesh1, [{"x"}, {}]>}, --// CHECK-SAME: %arg1: tensor<8x8xf32> {sdy.sharding = #sdy.sharding<@mesh2, [{}, {}]>}) --func.func @no_inlined_meshes_with_duplicates( -- %arg0: tensor<8x8xf32> {sdy.sharding = #sdy.sharding<@mesh1, [{"x"}, {}]>}, -- %arg1: tensor<8x8xf32> {sdy.sharding = #sdy.sharding<@mesh2, [{}, {}]>}) -> tensor<8x8xf32> { -- // CHECK-NEXT: stablehlo.add %arg0, %arg1 {sdy.sharding = #sdy.sharding_per_value<[<@mesh1, [{}, {"y"}]>]> -- %0 = stablehlo.add %arg0, %arg1 {sdy.sharding = #sdy.sharding_per_value<[<@mesh3, [{}, {"y"}]>]>} : tensor<8x8xf32> -- return %0 : tensor<8x8xf32> --} -- --// ----- -- --// CHECK: sdy.mesh @mesh = <["x"=2, "y"=4]> --sdy.mesh @mesh = <["x"=2, "y"=4]> -- --// CHECK: sdy.mesh @mesh_0 = <["x"=2, "y"=2]> -- --// CHECK-LABEL: func @inlined_mesh --func.func @inlined_mesh(%arg0: tensor<8x8xf32> {sdy.sharding = #sdy.sharding<@mesh, [{"x"}, {}]>}, %arg1: tensor<8x8xf32>) -> tensor<8x8xf32> { -- // CHECK-NEXT: stablehlo.add %arg0, %arg1 {sdy.sharding = #sdy.sharding_per_value<[<@mesh_0, [{"x"}, {}]>]> -- %0 = stablehlo.add %arg0, %arg1 {sdy.sharding = #sdy.sharding_per_value<[, [{"x"}, {}]>]>} : tensor<8x8xf32> -- return %0 : tensor<8x8xf32> --} -- --// ----- -- --// CHECK: sdy.mesh @mesh = <["x"=2, "y"=4]> --sdy.mesh @mesh = <["x"=2, "y"=4]> -- --// CHECK: sdy.mesh @mesh_0 = <["x"=2], device_ids=[1, 0]> --sdy.mesh @mesh_0 = <["x"=2], device_ids=[1, 0]> -- --// CHECK-NOT: @mesh_1 --sdy.mesh @mesh_1 = <["x"=2, "y"=4]> -- --// CHECK: sdy.mesh @maximal_mesh_5 = <[], device_ids=[5]> --sdy.mesh @maximal_mesh_5 = <[], device_ids=[5]> -- --// CHECK-NOT: @copy_of_mesh --sdy.mesh @copy_of_mesh = <["x"=2, "y"=4]> -- --// CHECK-NOT: @copy_of_mesh_0 --sdy.mesh @copy_of_mesh_0 = <["x"=2], device_ids=[1, 0]> -- --// CHECK: sdy.mesh @mesh_1 = <["a"=4]> --// CHECK: sdy.mesh @mesh_2 = <["x"=8]> --// CHECK: sdy.mesh @mesh_3 = <[], device_ids=[2]> -- --// CHECK-LABEL: func @many_inlined_meshes_and_duplicates( --// CHECK-SAME: %arg0: tensor<8x8xf32> {sdy.sharding = #sdy.sharding<@mesh_0, [{"x"}, {}]>}, --// CHECK-SAME: %arg1: tensor<8x8xf32> {sdy.sharding = #sdy.sharding<@mesh, [{"x"}, {?}], replicated={"y"}>}) --// CHECK-SAME: -> (tensor<8x8xf32> {sdy.sharding = #sdy.sharding<@mesh_2, [{?}, {?}]>}, --// CHECK-SAME: tensor<8x8xf32> {sdy.sharding = #sdy.sharding<@mesh_1, [{}, {}]>}) { --func.func @many_inlined_meshes_and_duplicates( -- %arg0: tensor<8x8xf32> {sdy.sharding = #sdy.sharding, [{"x"}, {}]>}, -- %arg1: tensor<8x8xf32> {sdy.sharding = #sdy.sharding<@mesh, [{"x"}, {?}], replicated={"y"}>}) -- -> (tensor<8x8xf32> {sdy.sharding = #sdy.sharding, [{?}, {?}]>}, -- tensor<8x8xf32> {sdy.sharding = #sdy.sharding, [{}, {}]>}) { -- // CHECK-NEXT: %[[ADD_0:.*]] = stablehlo.add %arg0, %arg0 {sdy.sharding = #sdy.sharding_per_value<[<@mesh, [{"x"}, {}]>]>} -- // CHECK-NEXT: %[[ADD_1:.*]] = stablehlo.add %[[ADD_0]], %arg1 : tensor<8x8xf32> -- // CHECK-NEXT: %[[SC_0:.*]] = sdy.sharding_constraint %[[ADD_1]] <@mesh, [{}, {"y"}]> -- // CHECK-NEXT: %[[ADD_2:.*]] = stablehlo.add %[[SC_0]], %[[SC_0]] {sdy.sharding = #sdy.sharding_per_value<[<@maximal_mesh_5, [{}, {}]>]>} -- // CHECK-NEXT: %[[SC_1:.*]] = sdy.sharding_constraint %[[ADD_2]] <@mesh_1, [{}, {"a"}]> -- // CHECK-NEXT: %[[ADD_3:.*]] = stablehlo.add %[[ADD_2]], %[[SC_1]] : tensor<8x8xf32> -- // CHECK-NEXT: return %[[SC_1]], %[[ADD_3]] -- %0 = stablehlo.add %arg0, %arg0 {sdy.sharding = #sdy.sharding_per_value<[<@mesh_1, [{"x"}, {}]>]>} : tensor<8x8xf32> -- %1 = stablehlo.add %0, %arg1 : tensor<8x8xf32> -- %2 = sdy.sharding_constraint %1 <@copy_of_mesh, [{}, {"y"}]> : tensor<8x8xf32> -- %3 = stablehlo.add %2, %2 {sdy.sharding = #sdy.sharding_per_value<[, [{}, {}]>]>} : tensor<8x8xf32> -- %4 = sdy.sharding_constraint %3 , [{}, {"a"}]> : tensor<8x8xf32> -- %5 = stablehlo.add %3, %4 : tensor<8x8xf32> -- return %4, %5 : tensor<8x8xf32>, tensor<8x8xf32> --} -- --// CHECK-LABEL: func @another_func_in_module( --// CHECK-SAME: %arg0: tensor<8x8xf32> {sdy.sharding = #sdy.sharding<@mesh_0, [{}, {}]>}) --// CHECK-SAME: -> (tensor<8x8xf32> {sdy.sharding = #sdy.sharding<@mesh_0, [{}, {}]>}) { --func.func @another_func_in_module( -- %arg0: tensor<8x8xf32> {sdy.sharding = #sdy.sharding<@mesh_0, [{}, {}]>}) -- -> (tensor<8x8xf32> {sdy.sharding = #sdy.sharding<@copy_of_mesh_0, [{}, {}]>}) { -- // CHECK-NEXT: %[[ADD:.*]] = stablehlo.add %arg0, %arg0 {sdy.sharding = #sdy.sharding_per_value<[<@mesh_1, [{"a"}, {}]>]>} -- // CHECK-NEXT: %[[SC:.*]] = sdy.sharding_constraint %[[ADD]] <@mesh_3, [{}, {}]> -- // CHECK-NEXT: return %[[SC]] -- %0 = stablehlo.add %arg0, %arg0 {sdy.sharding = #sdy.sharding_per_value<[, [{"a"}, {}]>]>} : tensor<8x8xf32> -- %1 = sdy.sharding_constraint %0 , [{}, {}]> : tensor<8x8xf32> -- return %1 : tensor<8x8xf32> --} -- --// ----- -- --// CHECK: sdy.mesh @mesh = <["x"=2]> --// CHECK: sdy.mesh @mesh_0 = <["y"=4]> --// CHECK: sdy.mesh @mesh_1 = <["x"=2], device_ids=[1, 0]> -- --// CHECK-LABEL: func @single_sharding_sdy_ops --func.func @single_sharding_sdy_ops(%arg0: tensor<8x8xf32>) -> tensor<8x8xf32> { -- // CHECK-NEXT: %[[SC:.*]] = sdy.data_flow_edge %arg0 sharding=<@mesh, [{"x"}, {}]> -- // CHECK-NEXT: %[[RESHARD:.*]] = sdy.reshard %[[SC]] <@mesh_0, [{}, {"y"}]> -- // CHECK-NEXT: sdy.sharding_constraint %[[RESHARD]] <@mesh_1, [{}, {"x"}]> -- %0 = sdy.data_flow_edge %arg0 sharding=, [{"x"}, {}]> : tensor<8x8xf32> -- %1 = sdy.reshard %0 , [{}, {"y"}]> : tensor<8x8xf32> -- %2 = sdy.sharding_constraint %1 , [{}, {"x"}]> : tensor<8x8xf32> -- return %2 : tensor<8x8xf32> --} -- --// ----- -- --// CHECK: sdy.mesh @mesh = <["x"=2]> -- --// CHECK-LABEL: func @manual_computation --func.func @manual_computation(%arg0: tensor<16x32xf32>, %arg1: tensor<16x32xf32>) -> tensor<16x32xf32> { -- // CHECK-NEXT: sdy.manual_computation(%arg0, %arg1) -- // CHECK-SAME: in_shardings=[<@mesh, [{"x"}, {}]>, <@mesh, [{"x"}, {}]>] -- // CHECK-SAME: out_shardings=[<@mesh, [{"x"}, {}]>] manual_axes={"x"} -- %0 = sdy.manual_computation(%arg0, %arg1) -- in_shardings=[, [{"x"}, {}]>, , [{"x"}, {}]>] -- out_shardings=[, [{"x"}, {}]>] manual_axes={"x"} -- (%arg2: tensor<8x32xf32>, %arg3: tensor<8x32xf32>) { -- %1 = stablehlo.add %arg2, %arg3 : tensor<8x32xf32> -- sdy.return %1 : tensor<8x32xf32> -- } : (tensor<16x32xf32>, tensor<16x32xf32>) -> tensor<16x32xf32> -- func.return %0: tensor<16x32xf32> --} -- --// ----- -- --// CHECK: sdy.mesh @mesh = <["x"=4]> --sdy.mesh @mesh = <["x"=4]> -- --// CHECK: sdy.mesh @mesh_0 = <["y"=8]> -- --// CHECK-LABEL: func @named_computation --func.func @named_computation(%arg0: tensor<16x32xf32>, %arg1: tensor<16x32xf32>) -> tensor<16x32xf32> { -- // CHECK-NEXT: sdy.named_computation<"foo">(%arg0, %arg1) -- // CHECK-SAME: in_shardings=[<@mesh, [{"x"}, {}]>, <@mesh, [{?}, {?}]>] -- // CHECK-SAME: out_shardings=[<@mesh_0, [{}, {"y"}]>] -- %0 = sdy.named_computation<"foo">(%arg0, %arg1) -- in_shardings=[<@mesh, [{"x"}, {}]>, , [{?}, {?}]>] -- out_shardings=[, [{}, {"y"}]>] -- (%arg2: tensor<16x32xf32>, %arg3: tensor<16x32xf32>) { -- %1 = stablehlo.add %arg2, %arg3 : tensor<16x32xf32> -- sdy.return %1 : tensor<16x32xf32> -- } : (tensor<16x32xf32>, tensor<16x32xf32>) -> tensor<16x32xf32> -- func.return %0: tensor<16x32xf32> --} -diff --git a/shardy/dialect/sdy/transforms/import/test/manual_axes_cleanup.mlir b/shardy/dialect/sdy/transforms/import/test/manual_axes_cleanup.mlir -index 81b8525..e97458a 100644 ---- a/shardy/dialect/sdy/transforms/import/test/manual_axes_cleanup.mlir -+++ b/shardy/dialect/sdy/transforms/import/test/manual_axes_cleanup.mlir -@@ -1,4 +1,4 @@ --// RUN: sdy_opt %s -sdy-manual-axes-cleanup | FileCheck %s -+// RUN: sdy_opt -split-input-file %s -sdy-manual-axes-cleanup | FileCheck %s - - sdy.mesh @mesh = <["c"=2, "a"=2, "b"=2]> - -@@ -6,7 +6,7 @@ sdy.mesh @mesh = <["c"=2, "a"=2, "b"=2]> - func.func @add_new_replicated(%arg0: tensor<8xf32>) -> tensor<8xf32> { - // CHECK-NEXT: sdy.manual_computation(%arg0) - // CHECK-SAME{LITERAL}: in_shardings=[<@mesh, [{"c", ?}], replicated={"a"}>] -- // CHECK-SAME{LITERAL}: out_shardings=[<@mesh, [{"c", ?}], replicated={"a"}>] -+ // CHECK-SAME{LITERAL}: out_shardings=[<@mesh, [{"c", ?}], replicated={"a"}>] - // CHECK-SAME{LITERAL}: manual_axes={"c", "a"} (%arg1: tensor<4xf32>) { - %0 = sdy.manual_computation(%arg0) in_shardings=[<@mesh, [{"c", ?}]>] out_shardings=[<@mesh, [{"c", ?}]>] manual_axes={"c", "a"} (%arg1: tensor<4xf32>) { - sdy.return %arg1 : tensor<4xf32> -@@ -19,7 +19,7 @@ func.func @add_new_replicated(%arg0: tensor<8xf32>) -> tensor<8xf32> { - func.func @add_to_existing_replicated(%arg0: tensor<8xf32>) -> tensor<8xf32> { - // CHECK-NEXT: sdy.manual_computation(%arg0) - // CHECK-SAME{LITERAL}: in_shardings=[<@mesh, [{"c", ?}], replicated={"a", "b"}>] -- // CHECK-SAME{LITERAL}: out_shardings=[<@mesh, [{"c", ?}], replicated={"a", "b"}>] -+ // CHECK-SAME{LITERAL}: out_shardings=[<@mesh, [{"c", ?}], replicated={"a", "b"}>] - // CHECK-SAME{LITERAL}: manual_axes={"c", "a", "b"} (%arg1: tensor<4xf32>) { - %0 = sdy.manual_computation(%arg0) in_shardings=[<@mesh, [{"c", ?}], replicated={"a"}>] out_shardings=[<@mesh, [{"c", ?}], replicated={"a"}>] manual_axes={"c", "a", "b"} (%arg1: tensor<4xf32>) { - sdy.return %arg1 : tensor<4xf32> -@@ -32,7 +32,7 @@ func.func @add_to_existing_replicated(%arg0: tensor<8xf32>) -> tensor<8xf32> { - func.func @add_to_existing_replicated_requires_sorting(%arg0: tensor<8xf32>) -> tensor<8xf32> { - // CHECK-NEXT: sdy.manual_computation(%arg0) - // CHECK-SAME{LITERAL}: in_shardings=[<@mesh, [{"c", ?}], replicated={"a", "b"}>] -- // CHECK-SAME{LITERAL}: out_shardings=[<@mesh, [{"c", ?}], replicated={"a", "b"}>] -+ // CHECK-SAME{LITERAL}: out_shardings=[<@mesh, [{"c", ?}], replicated={"a", "b"}>] - // CHECK-SAME{LITERAL}: manual_axes={"c", "a", "b"} (%arg1: tensor<4xf32>) { - %0 = sdy.manual_computation(%arg0) in_shardings=[<@mesh, [{"c", ?}], replicated={"b"}>] out_shardings=[<@mesh, [{"c", ?}], replicated={"a"}>] manual_axes={"c", "a", "b"} (%arg1: tensor<4xf32>) { - sdy.return %arg1 : tensor<4xf32> -@@ -46,27 +46,10 @@ func.func @add_to_existing_replicated_requires_sorting(%arg0: tensor<8xf32>) -> - func.func @sort_manual_axes_and_add_out_shardings(%arg0: tensor<8xf32>) -> tensor<8xf32> { - // CHECK-NEXT: sdy.manual_computation(%arg0) - // CHECK-SAME{LITERAL}: in_shardings=[<@mesh, [{"c", ?}], replicated={"a", "b"}>] -- // CHECK-SAME{LITERAL}: out_shardings=[<@mesh, [{"c", ?}], replicated={"a", "b"}>] -+ // CHECK-SAME{LITERAL}: out_shardings=[<@mesh, [{"c", ?}], replicated={"a", "b"}>] - // CHECK-SAME{LITERAL}: manual_axes={"c", "a", "b"} (%arg1: tensor<4xf32>) { - %0 = sdy.manual_computation(%arg0) in_shardings=[<@mesh, [{"c", ?}], replicated={"a", "b"}>] out_shardings=[<@mesh, [{"c", ?}], replicated={"b"}>] manual_axes={"b", "a", "c"} (%arg1: tensor<4xf32>) { - sdy.return %arg1 : tensor<4xf32> - } : (tensor<8xf32>) -> tensor<8xf32> - return %0 : tensor<8xf32> - } -- --sdy.mesh @mesh_xyz = <["z"=2, "x"=2, "y"=2]> -- --// CHECK-LABEL: @inlined_mesh_add_replicated_and_sort --func.func @inlined_mesh_add_replicated_and_sort(%arg0: tensor<8xf32>) -> tensor<8xf32> { -- // CHECK-NEXT: sdy.manual_computation(%arg0) -- // CHECK-SAME{LITERAL}: in_shardings=[<@mesh_xyz, [{"x", ?}], replicated={"z", "y"}>] -- // CHECK-SAME{LITERAL}: out_shardings=[, [{"z", ?}], replicated={"x", "y"}>] -- // CHECK-SAME{LITERAL}: manual_axes={"z", "x", "y"} (%arg1: tensor<4xf32>) { -- %0 = sdy.manual_computation(%arg0) -- in_shardings=[<@mesh_xyz, [{"x", ?}]>] -- out_shardings=[, [{"z", ?}], replicated={"y"}>] -- manual_axes={"y", "x", "z"} (%arg1: tensor<4xf32>) { -- sdy.return %arg1 : tensor<4xf32> -- } : (tensor<8xf32>) -> tensor<8xf32> -- return %0 : tensor<8xf32> --} -diff --git a/shardy/dialect/sdy/transforms/propagation/basic_propagation.cc b/shardy/dialect/sdy/transforms/propagation/basic_propagation.cc -index ff4f110..98043c4 100644 ---- a/shardy/dialect/sdy/transforms/propagation/basic_propagation.cc -+++ b/shardy/dialect/sdy/transforms/propagation/basic_propagation.cc -@@ -341,12 +341,12 @@ LogicalResult propagateFuncResults(FuncOp funcOp, - // op of `returnValue`. - // As such, we pass `returnValue` as both the operand and result. - returnValue, returnValue, getSharding(returnValue), -- getFuncResultSharding(funcOp, resNum), -+ funcOp.getResultAttrOfType(resNum, kShardingAttr), - [&](TensorShardingAttr sharding) { - setSharding(returnValue, sharding); - }, - [&](TensorShardingAttr sharding) { -- setFuncResultSharding(funcOp, resNum, sharding); -+ funcOp.setResultAttr(resNum, kShardingAttr, sharding); - }, - // Treat the sharding data flow b/w the `funcOp` terminator and func - // result attrs as an identity op. Create an equivalent sharding -diff --git a/shardy/dialect/sdy/transforms/propagation/test/propagation_pipeline.mlir b/shardy/dialect/sdy/transforms/propagation/test/propagation_pipeline.mlir -index fb26d23..4a51bd9 100644 ---- a/shardy/dialect/sdy/transforms/propagation/test/propagation_pipeline.mlir -+++ b/shardy/dialect/sdy/transforms/propagation/test/propagation_pipeline.mlir -@@ -1,4 +1,4 @@ --// RUN: sdy_opt %s -split-input-file -sdy-propagation-pipeline 2>&1 | FileCheck %s -+// RUN: sdy_opt %s -sdy-propagation-pipeline 2>&1 | FileCheck %s - - sdy.mesh @mesh = <["a"=2, "b"=2, "c"=2]> - -@@ -56,25 +56,3 @@ func.func @sharding_constraint_replaced_with_reshard(%arg0: tensor<8x8xf32>) -> - // CHECK-NEXT: return %arg0, %0 - return %arg0, %0 : tensor<8x8xf32>, tensor<8x8xf32> - } -- --// ----- -- --// CHECK: sdy.mesh @mesh = <["a"=2, "b"=2]> -- --// CHECK-LABEL: func @inlined_mesh( --// CHECK-SAME: %arg0: tensor<8x8xf32> {sdy.sharding = #sdy.sharding<@mesh, [{"a"}, {"b"}]>}, --// CHECK-SAME: %arg1: tensor<8x8xf32> {sdy.sharding = #sdy.sharding<@mesh, [{"a", ?}, {"b", ?}]>}, --// CHECK-SAME: %arg2: tensor<8x16xf32> {sdy.sharding = #sdy.sharding<@mesh, [{"b", ?}, {?}]>}) --// CHECK-SAME: -> (tensor<8x16xf32> {sdy.sharding = #sdy.sharding<@mesh, [{"a", ?}, {?}]>}) { --func.func @inlined_mesh( -- %arg0: tensor<8x8xf32> {sdy.sharding = #sdy.sharding, [{"a"}, {"b"}]>}, -- %arg1: tensor<8x8xf32>, %arg2: tensor<8x16xf32>) -> tensor<8x16xf32> { -- // CHECK-NEXT: %[[ADD:.*]] = stablehlo.add %arg0, %arg1 -- // CHECK-SAME: {sdy.sharding = #sdy.sharding_per_value<[<@mesh, [{"a", ?}, {"b", ?}]>]>} -- // CHECK-NEXT: stablehlo.dot_general %[[ADD]], %arg2 -- // CHECK-SAME: {sdy.sharding = #sdy.sharding_per_value<[<@mesh, [{"a", ?}, {?}]>]>} -- %0 = stablehlo.add %arg0, %arg1 : tensor<8x8xf32> -- %1 = stablehlo.dot_general %0, %arg2, contracting_dims = [1] x [0] : -- (tensor<8x8xf32>, tensor<8x16xf32>) -> tensor<8x16xf32> -- return %1 : tensor<8x16xf32> --} -diff --git a/shardy/dialect/sdy/transforms/propagation/user_priority_propagation.cc b/shardy/dialect/sdy/transforms/propagation/user_priority_propagation.cc -index 0cb1cfa..980cd8e 100644 ---- a/shardy/dialect/sdy/transforms/propagation/user_priority_propagation.cc -+++ b/shardy/dialect/sdy/transforms/propagation/user_priority_propagation.cc -@@ -24,6 +24,7 @@ limitations under the License. - #include "llvm/ADT/DenseSet.h" - #include "llvm/ADT/MapVector.h" - #include "llvm/ADT/STLExtras.h" -+#include "llvm/ADT/StringRef.h" - #include "llvm/Support/FormatVariadic.h" - #include "mlir/Dialect/Func/IR/FuncOps.h" - #include "mlir/IR/BuiltinAttributes.h" -@@ -36,6 +37,7 @@ limitations under the License. - #include "mlir/Support/LLVM.h" - #include "mlir/Support/LogicalResult.h" - #include "shardy/common/file_utils.h" -+#include "shardy/dialect/sdy/ir/constants.h" - #include "shardy/dialect/sdy/ir/dialect.h" - #include "shardy/dialect/sdy/ir/utils.h" - #include "shardy/dialect/sdy/transforms/propagation/auto_partitioner_registry.h" -@@ -117,7 +119,7 @@ TensorShardingAttr getUpdatedShardingForPriority( - return !seenAxesWithCurrentPriority.contains(axis); - }); - return TensorShardingAttr::get(curSharding.getContext(), -- curSharding.getMeshOrRef(), newDimShardings, -+ curSharding.getMeshName(), newDimShardings, - newReplicatedAxes); - } - -@@ -131,13 +133,14 @@ void updateReferencedShardingsForPriority( - getSharding(value), originalSharding, curPriority)); - } - -- for (const auto& [funcOp, resultShardings] : -+ for (auto [funcOp, resultShardings] : - shardingReferences.funcOpAndResultShardingsVec) { - for (auto [resNum, originalSharding] : resultShardings) { -- setFuncResultSharding( -- funcOp, resNum, -- getUpdatedShardingForPriority(getFuncResultSharding(funcOp, resNum), -- originalSharding, curPriority)); -+ funcOp.setResultAttr(resNum, kShardingAttr, -+ getUpdatedShardingForPriority( -+ funcOp.getResultAttrOfType( -+ resNum, kShardingAttr), -+ originalSharding, curPriority)); - } - } - } -@@ -165,14 +168,15 @@ TensorShardingAttr getInitializedSharding(TensorShardingAttr originalSharding, - dimSharding = DimensionShardingAttr::get(ctx, {}, /*isClosed=*/true); - } - } -- MeshAttr mesh = originalSharding.getMesh(op); -+ StringRef meshName = originalSharding.getMeshName(); -+ MeshAttr mesh = getMeshAttr(op, meshName); - assert(mesh && "unknown mesh"); - llvm::sort(newReplicatedAxes, AxisRefAttr::getMeshComparator(mesh)); - // TODO(tomnatan): we need to merge split axes and split them again when - // updating? or can we assume we won't see split axes? - -- return TensorShardingAttr::get(ctx, originalSharding.getMeshOrRef(), -- newDimShardings, newReplicatedAxes); -+ return TensorShardingAttr::get(ctx, meshName, newDimShardings, -+ newReplicatedAxes); - } - - // Clears `priorities` and add all non-zero priorities in `sharding` to it. -@@ -230,7 +234,8 @@ void addFuncResultShardingToPriorityMapAndInitialize( - FuncOp funcOp, int resNum, - PriorityToShardingReferences& priorityToShardingReferences, - llvm::SmallDenseSet& prioritiesInSharding) { -- auto sharding = getFuncResultSharding(funcOp, resNum); -+ auto sharding = -+ funcOp.getResultAttrOfType(resNum, kShardingAttr); - if (!sharding) { - return; - } -@@ -241,8 +246,8 @@ void addFuncResultShardingToPriorityMapAndInitialize( - funcOp) - .emplace_back(resNum, sharding); - } -- setFuncResultSharding(funcOp, resNum, -- getInitializedSharding(sharding, funcOp)); -+ funcOp.setResultAttr(resNum, kShardingAttr, -+ getInitializedSharding(sharding, funcOp)); - } - - // Traverses `funcOp` and for each value or func result with a sharding: -@@ -331,8 +336,8 @@ LogicalResult UserPriorityPropagationPassImpl::propagate( - SmallVector shardingReferencesPerPriority = - getShardingReferencesPerPriorityAndInitialize(moduleOp); - // We first run the first iteration (priority 0): -- if (failed(OpPriorityPropagationPassImpl::propagate( -- moduleOp, getDirectionToPropagate))) { -+ if (failed(OpPriorityPropagationPassImpl::propagate(moduleOp, -+ getDirectionToPropagate))) { - return failure(); - } - saveModuleOpAfterPriority(moduleOp, dumpDirectory, 0); -diff --git a/shardy/integrations/c/attributes.cc b/shardy/integrations/c/attributes.cc -index 3656d02..532df02 100644 ---- a/shardy/integrations/c/attributes.cc -+++ b/shardy/integrations/c/attributes.cc -@@ -35,12 +35,6 @@ AttrTy unwrapAttr(MlirAttribute attr) { - return mlir::cast(unwrap(attr)); - } - --template --mlir::ArrayRef unwrapAttrs(const MlirAttribute* attrs, -- intptr_t nAttrs) { -- return mlir::ArrayRef(reinterpret_cast(attrs), nAttrs); --} -- - } // namespace - - //===----------------------------------------------------------------------===// -@@ -75,9 +69,10 @@ bool sdyAttributeIsAMeshAttr(MlirAttribute attr) { - MlirAttribute sdyMeshAttrGet(MlirContext ctx, intptr_t nAxes, - const MlirAttribute* axes, intptr_t nDeviceIds, - const int64_t* deviceIds) { -- return wrap(sdy::MeshAttr::get(unwrap(ctx), -- unwrapAttrs(axes, nAxes), -- mlir::ArrayRef(deviceIds, nDeviceIds))); -+ return wrap(sdy::MeshAttr::get( -+ unwrap(ctx), -+ mlir::ArrayRef(reinterpret_cast(axes), nAxes), -+ mlir::ArrayRef(deviceIds, nDeviceIds))); - } - - int64_t sdyMeshAttrGetDeviceIdsSize(MlirAttribute attr) { -@@ -156,8 +151,9 @@ MlirAttribute sdyDimensionShardingAttrGet(MlirContext ctx, intptr_t nAxes, - const MlirAttribute* axes, - bool isClosed, int64_t priority) { - return wrap(sdy::DimensionShardingAttr::get( -- unwrap(ctx), unwrapAttrs(axes, nAxes), isClosed, -- priority == -1 ? std::nullopt : std::make_optional(priority))); -+ unwrap(ctx), -+ mlir::ArrayRef(reinterpret_cast(axes), nAxes), -+ isClosed, priority == -1 ? std::nullopt : std::make_optional(priority))); - } - - intptr_t sdyDimensionShardingAttrGetAxesSize(MlirAttribute attr) { -@@ -188,19 +184,23 @@ bool sdyAttributeIsATensorShardingAttr(MlirAttribute attr) { - return mlir::isa(unwrap(attr)); - } - --MlirAttribute sdyTensorShardingAttrGet(MlirContext ctx, MlirAttribute meshOrRef, -+MlirAttribute sdyTensorShardingAttrGet(MlirContext ctx, MlirStringRef meshName, - intptr_t nDimShardings, - const MlirAttribute* dimShardings, - intptr_t nReplicatedAxes, - const MlirAttribute* replicatedAxes) { - return wrap(sdy::TensorShardingAttr::get( -- unwrap(ctx), unwrap(meshOrRef), -- unwrapAttrs(dimShardings, nDimShardings), -- unwrapAttrs(replicatedAxes, nReplicatedAxes))); -+ unwrap(ctx), unwrap(meshName), -+ mlir::ArrayRef( -+ reinterpret_cast(dimShardings), -+ nDimShardings), -+ mlir::ArrayRef(reinterpret_cast(replicatedAxes), -+ nReplicatedAxes))); - } - --MlirAttribute sdyTensorShardingAttrGetMeshOrRef(MlirAttribute attr) { -- return wrap(unwrapAttr(attr).getMeshOrRef()); -+MlirStringRef sdyTensorShardingAttrGetMeshName(MlirAttribute attr) { -+ return wrap( -+ unwrapAttr(attr).getMeshSymName().getValue()); - } - - intptr_t sdyTensorShardingAttrGetDimShardingsSize(MlirAttribute attr) { -@@ -236,7 +236,9 @@ MlirAttribute sdyTensorShardingPerValueAttrGet(MlirContext ctx, - const MlirAttribute* shardings) { - return wrap(sdy::TensorShardingPerValueAttr::get( - unwrap(ctx), -- unwrapAttrs(shardings, nShardings))); -+ mlir::ArrayRef( -+ reinterpret_cast(shardings), -+ nShardings))); - } - - intptr_t sdyTensorShardingPerValueAttrGetShardingsSize(MlirAttribute attr) { -@@ -285,7 +287,9 @@ bool sdyAttributeIsATensorMappingAttr(MlirAttribute attr) { - MlirAttribute sdyTensorMappingAttrGet(MlirContext ctx, intptr_t nMappings, - const MlirAttribute* mappings) { - return wrap(sdy::TensorMappingAttr::get( -- unwrap(ctx), unwrapAttrs(mappings, nMappings))); -+ unwrap(ctx), -+ mlir::ArrayRef(reinterpret_cast(mappings), -+ nMappings))); - } - - intptr_t sdyTensorMappingAttrGetRank(MlirAttribute attr) { -@@ -318,8 +322,12 @@ MlirAttribute sdyOpShardingRuleAttrGet(MlirContext ctx, intptr_t nFactorSizes, - bool isCustomRule) { - return wrap(sdy::OpShardingRuleAttr::get( - unwrap(ctx), mlir::ArrayRef(factorSizes, nFactorSizes), -- unwrapAttrs(operandMappings, nOperandMappings), -- unwrapAttrs(resultMappings, nResultMappings), -+ mlir::ArrayRef( -+ reinterpret_cast(operandMappings), -+ nOperandMappings), -+ mlir::ArrayRef( -+ reinterpret_cast(resultMappings), -+ nResultMappings), - isCustomRule)); - } - -@@ -366,8 +374,8 @@ bool sdyAttributeIsAManualAxesAttr(MlirAttribute attr) { - - MlirAttribute sdyManualAxesAttrGet( - MlirContext ctx, intptr_t nAxes, const MlirAttribute* axes) { -- return wrap(sdy::ManualAxesAttr::get( -- unwrap(ctx), unwrapAttrs(axes, nAxes))); -+ return wrap(sdy::ManualAxesAttr::get(unwrap(ctx), mlir::ArrayRef( -+ reinterpret_cast(axes), nAxes))); - } - - intptr_t sdyManualAxesAttrGetAxesSize(MlirAttribute attr) { -diff --git a/shardy/integrations/c/attributes.h b/shardy/integrations/c/attributes.h -index 3f5c174..b3e9d76 100644 ---- a/shardy/integrations/c/attributes.h -+++ b/shardy/integrations/c/attributes.h -@@ -123,12 +123,12 @@ sdyDimensionShardingAttrGetPriority(MlirAttribute attr); - MLIR_CAPI_EXPORTED bool sdyAttributeIsATensorShardingAttr(MlirAttribute attr); - - MLIR_CAPI_EXPORTED MlirAttribute sdyTensorShardingAttrGet( -- MlirContext ctx, MlirAttribute meshOrRef, intptr_t nDimShardings, -+ MlirContext ctx, MlirStringRef meshName, intptr_t nDimShardings, - const MlirAttribute* dimShardings, intptr_t nReplicatedAxes, - const MlirAttribute* replicatedAxes); - --MLIR_CAPI_EXPORTED MlirAttribute --sdyTensorShardingAttrGetMeshOrRef(MlirAttribute attr); -+MLIR_CAPI_EXPORTED MlirStringRef -+sdyTensorShardingAttrGetMeshName(MlirAttribute attr); - - MLIR_CAPI_EXPORTED intptr_t - sdyTensorShardingAttrGetDimShardingsSize(MlirAttribute attr); -diff --git a/shardy/integrations/python/ir/sdy_module.cc b/shardy/integrations/python/ir/sdy_module.cc -index 8a933c6..04d276a 100644 ---- a/shardy/integrations/python/ir/sdy_module.cc -+++ b/shardy/integrations/python/ir/sdy_module.cc -@@ -15,10 +15,8 @@ limitations under the License. - - #include - #include --#include - #include - --#include "mlir-c/BuiltinAttributes.h" - #include "mlir-c/IR.h" - #include "mlir-c/Support.h" - #include "mlir/Bindings/Python/PybindAdaptors.h" // IWYU pragma: keep -@@ -55,15 +53,6 @@ MlirStringRef toStringRef(const std::string& s) { - return mlirStringRefCreate(s.c_str(), s.size()); - } - --MlirAttribute toMeshOrRefAttr( -- MlirContext ctx, -- const std::variant& meshOrRef) { -- if (auto* meshName = std::get_if(&meshOrRef)) { -- return mlirFlatSymbolRefAttrGet(ctx, toStringRef(*meshName)); -- } -- return std::get(meshOrRef); --} -- - PYBIND11_MODULE(_sdy, m) { - m.doc() = "SDY main Python extension"; - -@@ -203,26 +192,25 @@ PYBIND11_MODULE(_sdy, m) { - m, "TensorShardingAttr", sdyAttributeIsATensorShardingAttr) - .def_classmethod( - "get", -- [](py::object cls, -- const std::variant& meshOrRef, -+ [](py::object cls, const std::string& meshName, - const std::vector& dimensionShardings, - const std::vector& replicatedAxes, - MlirContext ctx) { - return cls(sdyTensorShardingAttrGet( -- ctx, toMeshOrRefAttr(ctx, meshOrRef), dimensionShardings.size(), -+ ctx, toStringRef(meshName), dimensionShardings.size(), - dimensionShardings.data(), replicatedAxes.size(), - replicatedAxes.data())); - }, -- py::arg("cls"), py::arg("mesh_or_ref"), -- py::arg("dimension_shardings"), -+ py::arg("cls"), py::arg("mesh_name"), py::arg("dimension_shardings"), - py::arg("replicated_axes") = std::vector(), - py::arg("context") = py::none(), -- "Creates a TensorShardingAttr with either an inlined mesh or mesh " -- "name, dimension shardings, and replicated axes.") -- .def_property_readonly("mesh_or_ref", -- [](MlirAttribute self) { -- return sdyTensorShardingAttrGetMeshOrRef(self); -- }) -+ "Creates a TensorShardingAttr with the mesh name, dimension " -+ "shardings, and replicated axes.") -+ .def_property_readonly( -+ "mesh_name", -+ [](MlirAttribute self) { -+ return toPyString(sdyTensorShardingAttrGetMeshName(self)); -+ }) - .def_property_readonly("dimension_shardings", - [](MlirAttribute self) { - return propertyVector( ++++ b/shardy/dialect/sdy/transforms/common/sharding_walker.h +@@ -29,8 +29,8 @@ using ConsumeShardingFn = std::function; + using TransformShardingFn = + std::function; + +-// Walks the given `rootOp` in forward pre-order and applies `consumeFn` on +-// any `TensorShardingAttr` encountered. ++// Walks the given `rootOp` in forward pre-order and applies `callback` on any ++// `TensorShardingAttr` encountered. + // + // In addition, applies `consumeOpFn` on every encountered op, before consuming + // its shardings. +@@ -39,7 +39,7 @@ void walkShardings( + ConsumeOpFn consumeOpFn = [](Operation*) {}); + + // Walks the given `rootOp` in forward pre-order and replaces any +-// `TensorShardingAttr` encountered with the result of applying `transformFn` on ++// `TensorShardingAttr` encountered with the result of applying `callback` on + // it. + // + // In addition, applies `consumeOpFn` on every encountered op, before diff --git a/third_party/llvm/generated.patch b/third_party/llvm/generated.patch -index 6525f92..f8fa2d1 100644 +index f8fa2d1..943d396 100644 --- a/third_party/llvm/generated.patch +++ b/third_party/llvm/generated.patch -@@ -1,149 +1,332 @@ +@@ -1,360 +1,419 @@ Auto generated patch. Do not edit or delete it, even if empty. --diff -ruN --strip-trailing-cr a/llvm/lib/Transforms/Vectorize/SLPVectorizer.cpp b/llvm/lib/Transforms/Vectorize/SLPVectorizer.cpp ----- a/llvm/lib/Transforms/Vectorize/SLPVectorizer.cpp --+++ b/llvm/lib/Transforms/Vectorize/SLPVectorizer.cpp --@@ -9181,12 +9181,13 @@ -- for (unsigned Cnt : Slices) { -- ArrayRef Slice = VL.slice(Cnt, VF); -- // If any instruction is vectorized already - do not try again. --- if (const TreeEntry *SE = getTreeEntry(Slice.front()); --+ if (TreeEntry *SE = getTreeEntry(Slice.front()); -- SE || getTreeEntry(Slice.back())) { -- if (!SE) -- continue; -- if (VF != SE->getVectorFactor() || !SE->isSame(Slice)) -- continue; --+ SE->UserTreeIndices.emplace_back(&E, UINT_MAX); -- AddCombinedNode(SE->Idx, Cnt); -- continue; -- } --@@ -13439,7 +13440,12 @@ -- if (CommonMask[Idx] != PoisonMaskElem) -- CommonMask[Idx] = Idx; -- for (auto [E, Idx] : SubVectors) { --- Value *V = castToScalarTyElem(E->VectorizedValue); --+ Value *V = E->VectorizedValue; --+ if (V->getType()->isIntOrIntVectorTy()) --+ V = castToScalarTyElem(V, any_of(E->Scalars, [&](Value *V) { --+ return !isKnownNonNegative( --+ V, SimplifyQuery(*R.DL)); --+ })); -- Vec = Builder.CreateInsertVector(Vec->getType(), Vec, V, -- Builder.getInt64(Idx)); -- if (!CommonMask.empty()) { --diff -ruN --strip-trailing-cr a/llvm/test/Transforms/SLPVectorizer/X86/subvector-minbitwidth-unsigned-value.ll b/llvm/test/Transforms/SLPVectorizer/X86/subvector-minbitwidth-unsigned-value.ll ----- a/llvm/test/Transforms/SLPVectorizer/X86/subvector-minbitwidth-unsigned-value.ll --+++ b/llvm/test/Transforms/SLPVectorizer/X86/subvector-minbitwidth-unsigned-value.ll --@@ -0,0 +1,97 @@ --+; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: --version 5 --+; RUN: opt -S --passes=slp-vectorizer -mtriple=x86_64-unknown-linux-gnu < %s | FileCheck %s -+diff -ruN --strip-trailing-cr a/libcxx/include/__memory/array_cookie.h b/libcxx/include/__memory/array_cookie.h -+--- a/libcxx/include/__memory/array_cookie.h -++++ b/libcxx/include/__memory/array_cookie.h -+@@ -24,7 +24,7 @@ -+ _LIBCPP_BEGIN_NAMESPACE_STD +-diff -ruN --strip-trailing-cr a/libcxx/include/__memory/array_cookie.h b/libcxx/include/__memory/array_cookie.h +---- a/libcxx/include/__memory/array_cookie.h +-+++ b/libcxx/include/__memory/array_cookie.h +-@@ -24,7 +24,7 @@ +- _LIBCPP_BEGIN_NAMESPACE_STD ++diff -ruN --strip-trailing-cr a/clang/test/SemaCXX/pragma-pack-packed-2.cpp b/clang/test/SemaCXX/pragma-pack-packed-2.cpp ++--- a/clang/test/SemaCXX/pragma-pack-packed-2.cpp +++++ b/clang/test/SemaCXX/pragma-pack-packed-2.cpp ++@@ -1,6 +1,6 @@ ++ // RUN: %clang_cc1 -triple s390x-ibm-zos -fzos-extensions -fsyntax-only -verify %s ++ // RUN: %clang_cc1 -triple s390x-ibm-zos -fzos-extensions -fxl-pragma-pack -fsyntax-only -verify %s ++-// RUN: %clang -target s390x-ibm-zos -S -emit-llvm -Xclang -verify -fno-xl-pragma-pack %s +++// RUN: %clang -target s390x-ibm-zos -S -emit-llvm -Xclang -verify -fno-xl-pragma-pack -o %t.ll %s + +- // Trait representing whether a type requires an array cookie at the start of its allocation when +--// allocated as `new T[n]` and deallocated as `delete array`. +-+// allocated as `new T[n]` and deallocated as `delete[] array`. +- // +- // Under the Itanium C++ ABI [1], we know that an array cookie is available unless `T` is trivially +- // destructible and the call to `operator delete[]` is not a sized operator delete. Under ABIs other +-diff -ruN --strip-trailing-cr a/libcxx/include/__memory/unique_ptr.h b/libcxx/include/__memory/unique_ptr.h +---- a/libcxx/include/__memory/unique_ptr.h +-+++ b/libcxx/include/__memory/unique_ptr.h +-@@ -102,6 +102,12 @@ +- }; ++ #pragma pack(show) // expected-warning {{value of #pragma pack(show) == 8}} ++ #pragma pack(twobyte) ++diff -ruN --strip-trailing-cr a/lld/ELF/Arch/AArch64.cpp b/lld/ELF/Arch/AArch64.cpp ++--- a/lld/ELF/Arch/AArch64.cpp +++++ b/lld/ELF/Arch/AArch64.cpp ++@@ -1208,10 +1208,12 @@ ++ } ++ } + +- template +-+struct __is_default_deleter : false_type {}; +-+ +-+template +-+struct __is_default_deleter > : true_type {}; +-+ +-+template +- struct __unique_ptr_deleter_sfinae { +- static_assert(!is_reference<_Deleter>::value, "incorrect specialization"); +- typedef const _Deleter& __lval_ref_type; +-@@ -307,11 +313,16 @@ +- // 1. When an array cookie (see [1]) exists at the beginning of the array allocation, we are +- // able to reuse that cookie to extract the size of the array and perform bounds checking. +- // An array cookie is a size inserted at the beginning of the allocation by the compiler. +--// That size is inserted implicitly when doing `new T[n]` in some cases, and its purpose +--// is to allow the runtime to destroy the `n` array elements when doing `delete array`. +-+// That size is inserted implicitly when doing `new T[n]` in some cases (as of writing this +-+// exactly when the array elements are not trivially destructible), and its main purpose is +-+// to allow the runtime to destroy the `n` array elements when doing `delete[] array`. +- // When we are able to use array cookies, we reuse information already available in the +- // current runtime, so bounds checking does not require changing libc++'s ABI. +- // +-+// However, note that we cannot assume the presence of an array cookie when a custom deleter +-+// is used, because the unique_ptr could have been created from an allocation that wasn't +-+// obtained via `new T[n]` (since it may not be deleted with `delete[] arr`). +-+// +- // 2. When the "bounded unique_ptr" ABI configuration (controlled by `_LIBCPP_ABI_BOUNDED_UNIQUE_PTR`) +- // is enabled, we store the size of the allocation (when it is known) so we can check it when +- // indexing into the `unique_ptr`. That changes the layout of `std::unique_ptr`, which is +-@@ -328,7 +339,7 @@ +- // try to fall back to using an array cookie when available. +- // +- // Finally, note that when this ABI configuration is enabled, we have no choice but to always +--// make space for a size to be stored in the unique_ptr. Indeed, while we might want to avoid +-+// make space for the size to be stored in the unique_ptr. Indeed, while we might want to avoid +- // storing the size when an array cookie is available, knowing whether an array cookie is available +- // requires the type stored in the unique_ptr to be complete, while unique_ptr can normally +- // accommodate incomplete types. +-@@ -339,7 +350,9 @@ +- __unique_ptr_array_bounds_stateless() = default; +- _LIBCPP_HIDE_FROM_ABI _LIBCPP_CONSTEXPR explicit __unique_ptr_array_bounds_stateless(size_t) {} ++-void elf::setAArch64TargetInfo(Ctx &ctx) { +++TargetInfo *elf::getAArch64TargetInfo(Ctx &ctx) { ++ if ((ctx.arg.andFeatures & GNU_PROPERTY_AARCH64_FEATURE_1_BTI) || ++- ctx.arg.zPacPlt) ++- ctx.target.reset(new AArch64BtiPac(ctx)); ++- else ++- ctx.target.reset(new AArch64(ctx)); +++ ctx.arg.zPacPlt) { +++ static AArch64BtiPac t(ctx); +++ return &t; +++ } +++ static AArch64 t(ctx); +++ return &t; ++ } ++diff -ruN --strip-trailing-cr a/lld/ELF/Arch/AMDGPU.cpp b/lld/ELF/Arch/AMDGPU.cpp ++--- a/lld/ELF/Arch/AMDGPU.cpp +++++ b/lld/ELF/Arch/AMDGPU.cpp ++@@ -219,4 +219,7 @@ ++ } ++ } + -+ // Trait representing whether a type requires an array cookie at the start of its allocation when -+-// allocated as `new T[n]` and deallocated as `delete array`. -++// allocated as `new T[n]` and deallocated as `delete[] array`. -+ // -+ // Under the Itanium C++ ABI [1], we know that an array cookie is available unless `T` is trivially -+ // destructible and the call to `operator delete[]` is not a sized operator delete. Under ABIs other -+diff -ruN --strip-trailing-cr a/libcxx/include/__memory/unique_ptr.h b/libcxx/include/__memory/unique_ptr.h -+--- a/libcxx/include/__memory/unique_ptr.h -++++ b/libcxx/include/__memory/unique_ptr.h -+@@ -102,6 +102,12 @@ -+ }; ++-void elf::setAMDGPUTargetInfo(Ctx &ctx) { ctx.target.reset(new AMDGPU(ctx)); } +++TargetInfo *elf::getAMDGPUTargetInfo(Ctx &ctx) { +++ static AMDGPU target(ctx); +++ return ⌖ +++} ++diff -ruN --strip-trailing-cr a/lld/ELF/Arch/ARM.cpp b/lld/ELF/Arch/ARM.cpp ++--- a/lld/ELF/Arch/ARM.cpp +++++ b/lld/ELF/Arch/ARM.cpp ++@@ -1533,7 +1533,10 @@ ++ "': " + toString(std::move(e))); ++ } + -+ template -++struct __is_default_deleter : false_type {}; -++ -++template -++struct __is_default_deleter > : true_type {}; -++ -++template -+ struct __unique_ptr_deleter_sfinae { -+ static_assert(!is_reference<_Deleter>::value, "incorrect specialization"); -+ typedef const _Deleter& __lval_ref_type; -+@@ -307,11 +313,16 @@ -+ // 1. When an array cookie (see [1]) exists at the beginning of the array allocation, we are -+ // able to reuse that cookie to extract the size of the array and perform bounds checking. -+ // An array cookie is a size inserted at the beginning of the allocation by the compiler. -+-// That size is inserted implicitly when doing `new T[n]` in some cases, and its purpose -+-// is to allow the runtime to destroy the `n` array elements when doing `delete array`. -++// That size is inserted implicitly when doing `new T[n]` in some cases (as of writing this -++// exactly when the array elements are not trivially destructible), and its main purpose is -++// to allow the runtime to destroy the `n` array elements when doing `delete[] array`. -+ // When we are able to use array cookies, we reuse information already available in the -+ // current runtime, so bounds checking does not require changing libc++'s ABI. -+ // -++// However, note that we cannot assume the presence of an array cookie when a custom deleter -++// is used, because the unique_ptr could have been created from an allocation that wasn't -++// obtained via `new T[n]` (since it may not be deleted with `delete[] arr`). -++// -+ // 2. When the "bounded unique_ptr" ABI configuration (controlled by `_LIBCPP_ABI_BOUNDED_UNIQUE_PTR`) -+ // is enabled, we store the size of the allocation (when it is known) so we can check it when -+ // indexing into the `unique_ptr`. That changes the layout of `std::unique_ptr`, which is -+@@ -328,7 +339,7 @@ -+ // try to fall back to using an array cookie when available. -+ // -+ // Finally, note that when this ABI configuration is enabled, we have no choice but to always -+-// make space for a size to be stored in the unique_ptr. Indeed, while we might want to avoid -++// make space for the size to be stored in the unique_ptr. Indeed, while we might want to avoid -+ // storing the size when an array cookie is available, knowing whether an array cookie is available -+ // requires the type stored in the unique_ptr to be complete, while unique_ptr can normally -+ // accommodate incomplete types. -+@@ -339,7 +350,9 @@ -+ __unique_ptr_array_bounds_stateless() = default; -+ _LIBCPP_HIDE_FROM_ABI _LIBCPP_CONSTEXPR explicit __unique_ptr_array_bounds_stateless(size_t) {} ++-void elf::setARMTargetInfo(Ctx &ctx) { ctx.target.reset(new ARM(ctx)); } +++TargetInfo *elf::getARMTargetInfo(Ctx &ctx) { +++ static ARM target(ctx); +++ return ⌖ +++} + +-- template ::value, int> = 0> +-+ template ::value && __has_array_cookie<_Tp>::value, int> = 0> +- _LIBCPP_HIDE_FROM_ABI _LIBCPP_CONSTEXPR bool __in_bounds(_Tp* __ptr, size_t __index) const { +- // In constant expressions, we can't check the array cookie so we just pretend that the index +- // is in-bounds. The compiler catches invalid accesses anyway. +-@@ -349,7 +362,9 @@ +- return __index < __cookie; ++ template void elf::writeARMCmseImportLib(Ctx &); ++ template void elf::writeARMCmseImportLib(Ctx &); ++diff -ruN --strip-trailing-cr a/lld/ELF/Arch/AVR.cpp b/lld/ELF/Arch/AVR.cpp ++--- a/lld/ELF/Arch/AVR.cpp +++++ b/lld/ELF/Arch/AVR.cpp ++@@ -267,7 +267,10 @@ + } ++ } + +-- template ::value, int> = 0> +-+ template ::value || !__has_array_cookie<_Tp>::value, int> = 0> +- _LIBCPP_HIDE_FROM_ABI _LIBCPP_CONSTEXPR bool __in_bounds(_Tp*, size_t) const { +- return true; // If we don't have an array cookie, we assume the access is in-bounds ++-void elf::setAVRTargetInfo(Ctx &ctx) { ctx.target.reset(new AVR(ctx)); } +++TargetInfo *elf::getAVRTargetInfo(Ctx &ctx) { +++ static AVR target(ctx); +++ return ⌖ +++} + -+- template ::value, int> = 0> -++ template ::value && __has_array_cookie<_Tp>::value, int> = 0> -+ _LIBCPP_HIDE_FROM_ABI _LIBCPP_CONSTEXPR bool __in_bounds(_Tp* __ptr, size_t __index) const { -+ // In constant expressions, we can't check the array cookie so we just pretend that the index -+ // is in-bounds. The compiler catches invalid accesses anyway. -+@@ -349,7 +362,9 @@ -+ return __index < __cookie; -+ } ++ static uint32_t getEFlags(InputFile *file) { ++ return cast>(file)->getObj().getHeader().e_flags; ++diff -ruN --strip-trailing-cr a/lld/ELF/Arch/Hexagon.cpp b/lld/ELF/Arch/Hexagon.cpp ++--- a/lld/ELF/Arch/Hexagon.cpp +++++ b/lld/ELF/Arch/Hexagon.cpp ++@@ -404,4 +404,7 @@ + } +-@@ -365,7 +380,9 @@ +- _LIBCPP_HIDE_FROM_ABI _LIBCPP_CONSTEXPR explicit __unique_ptr_array_bounds_stored(size_t __size) : __size_(__size) {} ++ } + +- // Use the array cookie if there's one +-- template ::value, int> = 0> +-+ template ::value && __has_array_cookie<_Tp>::value, int> = 0> +- _LIBCPP_HIDE_FROM_ABI _LIBCPP_CONSTEXPR bool __in_bounds(_Tp* __ptr, size_t __index) const { +- if (__libcpp_is_constant_evaluated()) +- return true; +-@@ -374,7 +391,9 @@ ++-void elf::setHexagonTargetInfo(Ctx &ctx) { ctx.target.reset(new Hexagon(ctx)); } +++TargetInfo *elf::getHexagonTargetInfo(Ctx &ctx) { +++ static Hexagon target(ctx); +++ return ⌖ +++} ++diff -ruN --strip-trailing-cr a/lld/ELF/Arch/LoongArch.cpp b/lld/ELF/Arch/LoongArch.cpp ++--- a/lld/ELF/Arch/LoongArch.cpp +++++ b/lld/ELF/Arch/LoongArch.cpp ++@@ -893,6 +893,7 @@ + } ++ } + -+- template ::value, int> = 0> -++ template ::value || !__has_array_cookie<_Tp>::value, int> = 0> -+ _LIBCPP_HIDE_FROM_ABI _LIBCPP_CONSTEXPR bool __in_bounds(_Tp*, size_t) const { -+ return true; // If we don't have an array cookie, we assume the access is in-bounds ++-void elf::setLoongArchTargetInfo(Ctx &ctx) { ++- ctx.target.reset(new LoongArch(ctx)); +++TargetInfo *elf::getLoongArchTargetInfo(Ctx &ctx) { +++ static LoongArch target(ctx); +++ return ⌖ ++ } ++diff -ruN --strip-trailing-cr a/lld/ELF/Arch/Mips.cpp b/lld/ELF/Arch/Mips.cpp ++--- a/lld/ELF/Arch/Mips.cpp +++++ b/lld/ELF/Arch/Mips.cpp ++@@ -779,23 +779,23 @@ ++ return cast>(file)->getObj().getHeader().e_flags & EF_MIPS_PIC; ++ } + +- // Otherwise, fall back on the stored size (if any) +-- template ::value, int> = 0> +-+ template ::value || !__has_array_cookie<_Tp>::value, int> = 0> +- _LIBCPP_HIDE_FROM_ABI _LIBCPP_CONSTEXPR bool __in_bounds(_Tp*, size_t __index) const { +- return __index < __size_; ++-void elf::setMipsTargetInfo(Ctx &ctx) { +++TargetInfo *elf::getMipsTargetInfo(Ctx &ctx) { ++ switch (ctx.arg.ekind) { ++ case ELF32LEKind: { ++- ctx.target.reset(new MIPS(ctx)); ++- return; +++ static MIPS t(ctx); +++ return &t; + } +-@@ -562,7 +581,7 @@ ++ case ELF32BEKind: { ++- ctx.target.reset(new MIPS(ctx)); ++- return; +++ static MIPS t(ctx); +++ return &t; + } +- +- _LIBCPP_HIDE_FROM_ABI _LIBCPP_CONSTEXPR_SINCE_CXX23 __add_lvalue_reference_t<_Tp> operator[](size_t __i) const { +-- _LIBCPP_ASSERT_VALID_ELEMENT_ACCESS(__checker_.__in_bounds(std::__to_address(__ptr_), __i), +-+ _LIBCPP_ASSERT_VALID_ELEMENT_ACCESS(__checker_.__in_bounds(std::__to_address(__ptr_), __i), +- "unique_ptr::operator[](index): index out of range"); +- return __ptr_[__i]; ++ case ELF64LEKind: { ++- ctx.target.reset(new MIPS(ctx)); ++- return; +++ static MIPS t(ctx); +++ return &t; + } -+@@ -365,7 +380,9 @@ -+ _LIBCPP_HIDE_FROM_ABI _LIBCPP_CONSTEXPR explicit __unique_ptr_array_bounds_stored(size_t __size) : __size_(__size) {} -+ -+ // Use the array cookie if there's one -+- template ::value, int> = 0> -++ template ::value && __has_array_cookie<_Tp>::value, int> = 0> -+ _LIBCPP_HIDE_FROM_ABI _LIBCPP_CONSTEXPR bool __in_bounds(_Tp* __ptr, size_t __index) const { -+ if (__libcpp_is_constant_evaluated()) -+ return true; -+@@ -374,7 +391,9 @@ ++ case ELF64BEKind: { ++- ctx.target.reset(new MIPS(ctx)); ++- return; +++ static MIPS t(ctx); +++ return &t; + } ++ default: ++ llvm_unreachable("unsupported target"); ++diff -ruN --strip-trailing-cr a/lld/ELF/Arch/MSP430.cpp b/lld/ELF/Arch/MSP430.cpp ++--- a/lld/ELF/Arch/MSP430.cpp +++++ b/lld/ELF/Arch/MSP430.cpp ++@@ -88,4 +88,7 @@ + } +-diff -ruN --strip-trailing-cr a/libcxx/test/std/utilities/smartptr/unique.ptr/unique.ptr.class/unique.ptr.observers/assert.subscript.pass.cpp b/libcxx/test/std/utilities/smartptr/unique.ptr/unique.ptr.class/unique.ptr.observers/assert.subscript.pass.cpp +---- a/libcxx/test/std/utilities/smartptr/unique.ptr/unique.ptr.class/unique.ptr.observers/assert.subscript.pass.cpp +-+++ b/libcxx/test/std/utilities/smartptr/unique.ptr/unique.ptr.class/unique.ptr.observers/assert.subscript.pass.cpp +-@@ -48,31 +48,24 @@ ++ } + +- template +- void test() { +-- // For types with an array cookie, we can always detect OOB accesses. +-+ // For types with an array cookie, we can always detect OOB accesses. Note that reliance on an array +-+ // cookie is limited to the default deleter, since a unique_ptr with a custom deleter may not have +-+ // been allocated with `new T[n]`. +- { +-- // Check with the default deleter +- { +-- { +-- std::unique_ptr ptr(new WithCookie[5]); +-- TEST_LIBCPP_ASSERT_FAILURE(ptr[6], "unique_ptr::operator[](index): index out of range"); +-- } +-- { +-- std::unique_ptr ptr = std::make_unique(5); +-- TEST_LIBCPP_ASSERT_FAILURE(ptr[6], "unique_ptr::operator[](index): index out of range"); +-- } +--#if TEST_STD_VER >= 20 +-- { +-- std::unique_ptr ptr = std::make_unique_for_overwrite(5); +-- TEST_LIBCPP_ASSERT_FAILURE(ptr[6] = WithCookie(), "unique_ptr::operator[](index): index out of range"); +-- } +--#endif +-+ std::unique_ptr ptr(new WithCookie[5]); +-+ TEST_LIBCPP_ASSERT_FAILURE(ptr[6], "unique_ptr::operator[](index): index out of range"); +- } +-- +-- // Check with a custom deleter +- { +-- std::unique_ptr ptr(new WithCookie[5]); +-+ std::unique_ptr ptr = std::make_unique(5); +- TEST_LIBCPP_ASSERT_FAILURE(ptr[6], "unique_ptr::operator[](index): index out of range"); +- } +-+#if TEST_STD_VER >= 20 +-+ { +-+ std::unique_ptr ptr = std::make_unique_for_overwrite(5); +-+ TEST_LIBCPP_ASSERT_FAILURE(ptr[6] = WithCookie(), "unique_ptr::operator[](index): index out of range"); +-+ } +-+#endif ++-void elf::setMSP430TargetInfo(Ctx &ctx) { ctx.target.reset(new MSP430(ctx)); } +++TargetInfo *elf::getMSP430TargetInfo(Ctx &ctx) { +++ static MSP430 target(ctx); +++ return ⌖ +++} ++diff -ruN --strip-trailing-cr a/lld/ELF/Arch/PPC64.cpp b/lld/ELF/Arch/PPC64.cpp ++--- a/lld/ELF/Arch/PPC64.cpp +++++ b/lld/ELF/Arch/PPC64.cpp ++@@ -1747,4 +1747,7 @@ ++ return true; ++ } + -+ // Otherwise, fall back on the stored size (if any) -+- template ::value, int> = 0> -++ template ::value || !__has_array_cookie<_Tp>::value, int> = 0> -+ _LIBCPP_HIDE_FROM_ABI _LIBCPP_CONSTEXPR bool __in_bounds(_Tp*, size_t __index) const { -+ return __index < __size_; -+ } -+@@ -562,7 +581,7 @@ -+ } ++-void elf::setPPC64TargetInfo(Ctx &ctx) { ctx.target.reset(new PPC64(ctx)); } +++TargetInfo *elf::getPPC64TargetInfo(Ctx &ctx) { +++ static PPC64 target(ctx); +++ return ⌖ +++} ++diff -ruN --strip-trailing-cr a/lld/ELF/Arch/PPC.cpp b/lld/ELF/Arch/PPC.cpp ++--- a/lld/ELF/Arch/PPC.cpp +++++ b/lld/ELF/Arch/PPC.cpp ++@@ -523,4 +523,7 @@ + } ++ } + +- // For types that don't have an array cookie, things are a bit more complicated. We can detect OOB accesses +-@@ -97,14 +90,9 @@ +- #endif ++-void elf::setPPCTargetInfo(Ctx &ctx) { ctx.target.reset(new PPC(ctx)); } +++TargetInfo *elf::getPPCTargetInfo(Ctx &ctx) { +++ static PPC target(ctx); +++ return ⌖ +++} ++diff -ruN --strip-trailing-cr a/lld/ELF/Arch/RISCV.cpp b/lld/ELF/Arch/RISCV.cpp ++--- a/lld/ELF/Arch/RISCV.cpp +++++ b/lld/ELF/Arch/RISCV.cpp ++@@ -1329,4 +1329,7 @@ ++ mergeAttributesSection(ctx, sections)); ++ } + -+ _LIBCPP_HIDE_FROM_ABI _LIBCPP_CONSTEXPR_SINCE_CXX23 __add_lvalue_reference_t<_Tp> operator[](size_t __i) const { -+- _LIBCPP_ASSERT_VALID_ELEMENT_ACCESS(__checker_.__in_bounds(std::__to_address(__ptr_), __i), -++ _LIBCPP_ASSERT_VALID_ELEMENT_ACCESS(__checker_.__in_bounds(std::__to_address(__ptr_), __i), -+ "unique_ptr::operator[](index): index out of range"); -+ return __ptr_[__i]; ++-void elf::setRISCVTargetInfo(Ctx &ctx) { ctx.target.reset(new RISCV(ctx)); } +++TargetInfo *elf::getRISCVTargetInfo(Ctx &ctx) { +++ static RISCV target(ctx); +++ return ⌖ +++} ++diff -ruN --strip-trailing-cr a/lld/ELF/Arch/SPARCV9.cpp b/lld/ELF/Arch/SPARCV9.cpp ++--- a/lld/ELF/Arch/SPARCV9.cpp +++++ b/lld/ELF/Arch/SPARCV9.cpp ++@@ -193,4 +193,7 @@ ++ relocateNoSym(buf + 4, R_SPARC_WDISP19, -(off + 4 - pltEntrySize)); ++ } + +- // Make sure that we carry the bounds information properly through conversions, assignments, etc. +-- // These tests are mostly relevant when the ABI setting is enabled (with a stateful bounds-checker), +-- // but we still run them for types with an array cookie either way. +-+ // These tests are only relevant when the ABI setting is enabled (with a stateful bounds-checker). +- #if defined(_LIBCPP_ABI_BOUNDED_UNIQUE_PTR) +-- using Types = types::type_list; +--#else +-- using Types = types::type_list; +--#endif +-- types::for_each(Types(), [] { +-+ types::for_each(types::type_list(), [] { +- // Bounds carried through move construction +- { +- std::unique_ptr ptr = std::make_unique(5); +-@@ -135,6 +123,7 @@ +- TEST_LIBCPP_ASSERT_FAILURE(other[6], "unique_ptr::operator[](index): index out of range"); +- } +- }); +-+#endif ++-void elf::setSPARCV9TargetInfo(Ctx &ctx) { ctx.target.reset(new SPARCV9(ctx)); } +++TargetInfo *elf::getSPARCV9TargetInfo(Ctx &ctx) { +++ static SPARCV9 target(ctx); +++ return ⌖ +++} ++diff -ruN --strip-trailing-cr a/lld/ELF/Arch/SystemZ.cpp b/lld/ELF/Arch/SystemZ.cpp ++--- a/lld/ELF/Arch/SystemZ.cpp +++++ b/lld/ELF/Arch/SystemZ.cpp ++@@ -600,4 +600,7 @@ + } -+diff -ruN --strip-trailing-cr a/libcxx/test/std/utilities/smartptr/unique.ptr/unique.ptr.class/unique.ptr.observers/assert.subscript.pass.cpp b/libcxx/test/std/utilities/smartptr/unique.ptr/unique.ptr.class/unique.ptr.observers/assert.subscript.pass.cpp -+--- a/libcxx/test/std/utilities/smartptr/unique.ptr/unique.ptr.class/unique.ptr.observers/assert.subscript.pass.cpp -++++ b/libcxx/test/std/utilities/smartptr/unique.ptr/unique.ptr.class/unique.ptr.observers/assert.subscript.pass.cpp -+@@ -48,31 +48,24 @@ -+ -+ template -+ void test() { -+- // For types with an array cookie, we can always detect OOB accesses. -++ // For types with an array cookie, we can always detect OOB accesses. Note that reliance on an array -++ // cookie is limited to the default deleter, since a unique_ptr with a custom deleter may not have -++ // been allocated with `new T[n]`. -+ { -+- // Check with the default deleter -+ { -+- { -+- std::unique_ptr ptr(new WithCookie[5]); -+- TEST_LIBCPP_ASSERT_FAILURE(ptr[6], "unique_ptr::operator[](index): index out of range"); -+- } -+- { -+- std::unique_ptr ptr = std::make_unique(5); -+- TEST_LIBCPP_ASSERT_FAILURE(ptr[6], "unique_ptr::operator[](index): index out of range"); -+- } -+-#if TEST_STD_VER >= 20 -+- { -+- std::unique_ptr ptr = std::make_unique_for_overwrite(5); -+- TEST_LIBCPP_ASSERT_FAILURE(ptr[6] = WithCookie(), "unique_ptr::operator[](index): index out of range"); -+- } -+-#endif -++ std::unique_ptr ptr(new WithCookie[5]); -++ TEST_LIBCPP_ASSERT_FAILURE(ptr[6], "unique_ptr::operator[](index): index out of range"); -+ } -+- -+- // Check with a custom deleter -+ { -+- std::unique_ptr ptr(new WithCookie[5]); -++ std::unique_ptr ptr = std::make_unique(5); -+ TEST_LIBCPP_ASSERT_FAILURE(ptr[6], "unique_ptr::operator[](index): index out of range"); -+ } -++#if TEST_STD_VER >= 20 -++ { -++ std::unique_ptr ptr = std::make_unique_for_overwrite(5); -++ TEST_LIBCPP_ASSERT_FAILURE(ptr[6] = WithCookie(), "unique_ptr::operator[](index): index out of range"); + } + +- template +-diff -ruN --strip-trailing-cr a/libcxx/test/std/utilities/smartptr/unique.ptr/unique.ptr.class/unique.ptr.observers/op_subscript.runtime.pass.cpp b/libcxx/test/std/utilities/smartptr/unique.ptr/unique.ptr.class/unique.ptr.observers/op_subscript.runtime.pass.cpp +---- a/libcxx/test/std/utilities/smartptr/unique.ptr/unique.ptr.class/unique.ptr.observers/op_subscript.runtime.pass.cpp +-+++ b/libcxx/test/std/utilities/smartptr/unique.ptr/unique.ptr.class/unique.ptr.observers/op_subscript.runtime.pass.cpp +-@@ -46,6 +46,11 @@ +- template +- struct CustomDeleter : std::default_delete {}; ++-void elf::setSystemZTargetInfo(Ctx &ctx) { ctx.target.reset(new SystemZ(ctx)); } +++TargetInfo *elf::getSystemZTargetInfo(Ctx &ctx) { +++ static SystemZ t(ctx); +++ return &t; +++} ++diff -ruN --strip-trailing-cr a/lld/ELF/Arch/X86_64.cpp b/lld/ELF/Arch/X86_64.cpp ++--- a/lld/ELF/Arch/X86_64.cpp +++++ b/lld/ELF/Arch/X86_64.cpp ++@@ -1237,17 +1237,21 @@ ++ write32le(buf + 8, ctx.in.plt->getVA() - pltEntryAddr - 12); ++ } + +-+struct NoopDeleter { +-+ template +-+ TEST_CONSTEXPR_CXX23 void operator()(T*) const {} +-+}; +-+ +- TEST_CONSTEXPR_CXX23 bool test() { +- // Basic test +- { +-@@ -112,12 +117,33 @@ +- WithNonTrivialDtor<16>, +- WithNonTrivialDtor<256>>; +- types::for_each(TrickyCookieTypes(), [] { +-- types::for_each(types::type_list, CustomDeleter>(), [] { +-- std::unique_ptr p(new T[3]); +-+ // Array allocated with `new T[n]`, default deleter +-+ { +-+ std::unique_ptr> p(new T[3]); +-+ assert(p[0] == T()); +-+ assert(p[1] == T()); +-+ assert(p[2] == T()); +-+ } +-+ +-+ // Array allocated with `new T[n]`, custom deleter +-+ { +-+ std::unique_ptr> p(new T[3]); +-+ assert(p[0] == T()); +-+ assert(p[1] == T()); +-+ assert(p[2] == T()); +-+ } +-+ +-+ // Array not allocated with `new T[n]`, custom deleter +-+ // +-+ // This test aims to ensure that the implementation doesn't try to use an array cookie +-+ // when there is none. +-+ { +-+ T array[50] = {}; +-+ std::unique_ptr p(&array[0]); +- assert(p[0] == T()); +- assert(p[1] == T()); +- assert(p[2] == T()); +-- }); +-+ } +- }); ++-void elf::setX86_64TargetInfo(Ctx &ctx) { +++TargetInfo *elf::getX86_64TargetInfo(Ctx &ctx) { ++ if (ctx.arg.zRetpolineplt) { ++- if (ctx.arg.zNow) ++- ctx.target.reset(new RetpolineZNow(ctx)); ++- else ++- ctx.target.reset(new Retpoline(ctx)); ++- return; +++ if (ctx.arg.zNow) { +++ static RetpolineZNow t(ctx); +++ return &t; ++ } -++#endif +++ static Retpoline t(ctx); +++ return &t; + } +- #endif // C++20 +-diff -ruN --strip-trailing-cr a/lldb/test/API/functionalities/data-formatter/data-formatter-stl/libcxx-simulators/string/main.cpp b/lldb/test/API/functionalities/data-formatter/data-formatter-stl/libcxx-simulators/string/main.cpp +---- a/lldb/test/API/functionalities/data-formatter/data-formatter-stl/libcxx-simulators/string/main.cpp +-+++ b/lldb/test/API/functionalities/data-formatter/data-formatter-stl/libcxx-simulators/string/main.cpp +-@@ -20,11 +20,7 @@ +- // Pre-D128285 layout. +- #define PACKED_ANON_STRUCT +- #endif +--#if REVISION <= 4 +--// Pre-2a1ef74 layout. +--#define NON_STANDARD_PADDING +--#endif +--// REVISION == 5: current layout +-+// REVISION == 4: current layout + +- #ifdef PACKED_ANON_STRUCT +- #define BEGIN_PACKED_ANON_STRUCT struct __attribute__((packed)) { +-@@ -38,7 +34,6 @@ +- namespace std { +- namespace __lldb { ++- if (ctx.arg.andFeatures & GNU_PROPERTY_X86_FEATURE_1_IBT) ++- ctx.target.reset(new IntelIBT(ctx)); ++- else ++- ctx.target.reset(new X86_64(ctx)); +++ if (ctx.arg.andFeatures & GNU_PROPERTY_X86_FEATURE_1_IBT) { +++ static IntelIBT t(ctx); +++ return &t; +++ } +++ +++ static X86_64 t(ctx); +++ return &t; ++ } ++diff -ruN --strip-trailing-cr a/lld/ELF/Arch/X86.cpp b/lld/ELF/Arch/X86.cpp ++--- a/lld/ELF/Arch/X86.cpp +++++ b/lld/ELF/Arch/X86.cpp ++@@ -706,17 +706,21 @@ ++ write32le(buf + 22, -off - 26); ++ } + +--#ifdef NON_STANDARD_PADDING +- #if defined(ALTERNATE_LAYOUT) && defined(SUBCLASS_PADDING) +- template struct __padding { +- unsigned char __xx[sizeof(_CharT) - 1]; +-@@ -46,13 +41,6 @@ ++-void elf::setX86TargetInfo(Ctx &ctx) { +++TargetInfo *elf::getX86TargetInfo(Ctx &ctx) { ++ if (ctx.arg.zRetpolineplt) { ++- if (ctx.arg.isPic) ++- ctx.target.reset(new RetpolinePic(ctx)); ++- else ++- ctx.target.reset(new RetpolineNoPic(ctx)); ++- return; +++ if (ctx.arg.isPic) { +++ static RetpolinePic t(ctx); +++ return &t; +++ } +++ static RetpolineNoPic t(ctx); +++ return &t; + } -+ -+ // For types that don't have an array cookie, things are a bit more complicated. We can detect OOB accesses -+@@ -97,14 +90,9 @@ -+ #endif -+ -+ // Make sure that we carry the bounds information properly through conversions, assignments, etc. -+- // These tests are mostly relevant when the ABI setting is enabled (with a stateful bounds-checker), -+- // but we still run them for types with an array cookie either way. -++ // These tests are only relevant when the ABI setting is enabled (with a stateful bounds-checker). -+ #if defined(_LIBCPP_ABI_BOUNDED_UNIQUE_PTR) -+- using Types = types::type_list; -+-#else -+- using Types = types::type_list; -+-#endif -+- types::for_each(Types(), [] { -++ types::for_each(types::type_list(), [] { -+ // Bounds carried through move construction -+ { -+ std::unique_ptr ptr = std::make_unique(5); -+@@ -135,6 +123,7 @@ -+ TEST_LIBCPP_ASSERT_FAILURE(other[6], "unique_ptr::operator[](index): index out of range"); -+ } -+ }); -++#endif + +- template struct __padding<_CharT, 1> {}; +- #endif +--#else // !NON_STANDARD_PADDING +--template struct __padding { +-- char __padding_[_PaddingSize]; +--}; +-- +--template <> struct __padding<0> {}; +--#endif ++- if (ctx.arg.andFeatures & GNU_PROPERTY_X86_FEATURE_1_IBT) ++- ctx.target.reset(new IntelIBT(ctx)); ++- else ++- ctx.target.reset(new X86(ctx)); +++ if (ctx.arg.andFeatures & GNU_PROPERTY_X86_FEATURE_1_IBT) { +++ static IntelIBT t(ctx); +++ return &t; +++ } +++ +++ static X86 t(ctx); +++ return &t; + } ++diff -ruN --strip-trailing-cr a/lld/ELF/Config.h b/lld/ELF/Config.h ++--- a/lld/ELF/Config.h +++++ b/lld/ELF/Config.h ++@@ -545,7 +545,7 @@ ++ Config arg; ++ LinkerDriver driver; ++ LinkerScript *script; ++- std::unique_ptr target; +++ TargetInfo *target; + -+ template -+diff -ruN --strip-trailing-cr a/libcxx/test/std/utilities/smartptr/unique.ptr/unique.ptr.class/unique.ptr.observers/op_subscript.runtime.pass.cpp b/libcxx/test/std/utilities/smartptr/unique.ptr/unique.ptr.class/unique.ptr.observers/op_subscript.runtime.pass.cpp -+--- a/libcxx/test/std/utilities/smartptr/unique.ptr/unique.ptr.class/unique.ptr.observers/op_subscript.runtime.pass.cpp -++++ b/libcxx/test/std/utilities/smartptr/unique.ptr/unique.ptr.class/unique.ptr.observers/op_subscript.runtime.pass.cpp -+@@ -46,6 +46,11 @@ -+ template -+ struct CustomDeleter : std::default_delete {}; -+ -++struct NoopDeleter { -++ template -++ TEST_CONSTEXPR_CXX23 void operator()(T*) const {} -++}; - + --+define i1 @test(i64 %v1, ptr %v2, i32 %v3, i1 %v4) { --+; CHECK-LABEL: define i1 @test( --+; CHECK-SAME: i64 [[V1:%.*]], ptr [[V2:%.*]], i32 [[V3:%.*]], i1 [[V4:%.*]]) { --+; CHECK-NEXT: [[NEWFUNCROOT:.*:]] --+; CHECK-NEXT: [[TMP0:%.*]] = insertelement <2 x i64> poison, i64 [[V1]], i32 0 --+; CHECK-NEXT: [[TMP1:%.*]] = shufflevector <2 x i64> [[TMP0]], <2 x i64> poison, <2 x i32> zeroinitializer --+; CHECK-NEXT: [[TMP2:%.*]] = lshr <2 x i64> [[TMP1]], --+; CHECK-NEXT: [[TMP3:%.*]] = trunc <2 x i64> [[TMP2]] to <2 x i8> --+; CHECK-NEXT: [[TMP4:%.*]] = and <2 x i8> [[TMP3]], --+; CHECK-NEXT: [[TMP5:%.*]] = zext <2 x i8> [[TMP4]] to <2 x i32> --+; CHECK-NEXT: [[TMP6:%.*]] = icmp eq <2 x i32> [[TMP5]], zeroinitializer --+; CHECK-NEXT: [[TMP7:%.*]] = insertelement <4 x i32> poison, i32 [[V3]], i32 0 --+; CHECK-NEXT: [[TMP8:%.*]] = shufflevector <4 x i32> [[TMP7]], <4 x i32> poison, <4 x i32> --+; CHECK-NEXT: [[TMP9:%.*]] = zext <2 x i8> [[TMP4]] to <2 x i32> --+; CHECK-NEXT: [[TMP10:%.*]] = call <4 x i32> @llvm.vector.insert.v4i32.v2i32(<4 x i32> [[TMP8]], <2 x i32> [[TMP9]], i64 0) --+; CHECK-NEXT: [[TMP11:%.*]] = uitofp <4 x i32> [[TMP10]] to <4 x float> --+; CHECK-NEXT: [[TMP12:%.*]] = fdiv <4 x float> zeroinitializer, [[TMP11]] --+; CHECK-NEXT: [[TMP13:%.*]] = insertelement <4 x i1> poison, i1 [[V4]], i32 0 --+; CHECK-NEXT: [[TMP14:%.*]] = shufflevector <4 x i1> [[TMP13]], <4 x i1> poison, <4 x i32> --+; CHECK-NEXT: [[TMP15:%.*]] = call <4 x i1> @llvm.vector.insert.v4i1.v2i1(<4 x i1> [[TMP14]], <2 x i1> [[TMP6]], i64 0) --+; CHECK-NEXT: [[TMP16:%.*]] = select <4 x i1> [[TMP15]], <4 x float> zeroinitializer, <4 x float> [[TMP12]] --+; CHECK-NEXT: [[TMP17:%.*]] = extractelement <4 x float> [[TMP16]], i32 3 --+; CHECK-NEXT: [[CONV_I_I1743_3:%.*]] = fptoui float [[TMP17]] to i32 --+; CHECK-NEXT: [[TMP18:%.*]] = icmp ne i32 [[CONV_I_I1743_3]], 0 --+; CHECK-NEXT: [[TMP19:%.*]] = bitcast <4 x float> [[TMP16]] to <4 x i32> --+; CHECK-NEXT: [[TMP20:%.*]] = icmp ult <4 x i32> [[TMP19]], --+; CHECK-NEXT: [[TMP21:%.*]] = extractelement <4 x i1> [[TMP20]], i32 3 --+; CHECK-NEXT: [[NARROW:%.*]] = select i1 [[TMP21]], i1 [[TMP18]], i1 false --+; CHECK-NEXT: [[TMP22:%.*]] = extractelement <4 x float> [[TMP16]], i32 2 --+; CHECK-NEXT: [[CONV_I_I1743_2:%.*]] = fptoui float [[TMP22]] to i32 --+; CHECK-NEXT: [[TMP23:%.*]] = extractelement <4 x i1> [[TMP20]], i32 2 --+; CHECK-NEXT: [[NARROW1:%.*]] = select i1 [[TMP23]], i32 [[CONV_I_I1743_2]], i32 0 --+; CHECK-NEXT: [[TMP24:%.*]] = zext i1 [[NARROW]] to i32 --+; CHECK-NEXT: [[TMP25:%.*]] = or i32 [[NARROW1]], [[TMP24]] --+; CHECK-NEXT: [[TMP26:%.*]] = extractelement <4 x float> [[TMP16]], i32 1 --+; CHECK-NEXT: [[CONV_I_I1743_1:%.*]] = fptoui float [[TMP26]] to i32 --+; CHECK-NEXT: [[TMP27:%.*]] = extractelement <4 x i1> [[TMP20]], i32 1 --+; CHECK-NEXT: [[NARROW2:%.*]] = select i1 [[TMP27]], i32 [[CONV_I_I1743_1]], i32 0 --+; CHECK-NEXT: [[RV3:%.*]] = or i32 [[TMP25]], [[NARROW2]] --+; CHECK-NEXT: [[TMP28:%.*]] = extractelement <4 x float> [[TMP16]], i32 0 --+; CHECK-NEXT: [[CONV_I_I1743:%.*]] = fptoui float [[TMP28]] to i32 --+; CHECK-NEXT: [[TMP29:%.*]] = extractelement <4 x i1> [[TMP20]], i32 0 --+; CHECK-NEXT: [[NARROW4:%.*]] = select i1 [[TMP29]], i32 [[CONV_I_I1743]], i32 0 --+; CHECK-NEXT: [[RT5:%.*]] = or i32 [[RV3]], [[NARROW4]] --+; CHECK-NEXT: [[RT:%.*]] = zext i32 [[RT5]] to i64 --+; CHECK-NEXT: store i64 [[RT]], ptr [[V2]], align 1 --+; CHECK-NEXT: ret i1 false --+; --+newFuncRoot: --+ %conv.i147.i1756.3 = uitofp i32 %v3 to float --+ %div.i.i.i1749.3 = fdiv float 0.000000e+00, %conv.i147.i1756.3 --+ %cond.i.i.i1751.3 = select i1 %v4, float 0.000000e+00, float %div.i.i.i1749.3 --+ %conv.i147.i1756.2 = uitofp i32 %v3 to float --+ %div.i.i.i1749.2 = fdiv float 0.000000e+00, %conv.i147.i1756.2 --+ %cond.i.i.i1751.2 = select i1 %v4, float 0.000000e+00, float %div.i.i.i1749.2 --+ %0 = lshr i64 %v1, 40 --+ %1 = trunc i64 %0 to i32 --+ %tt2 = and i32 %1, 255 --+ %cmp1.i.i.i1746.1 = icmp eq i32 %tt2, 0 --+ %conv.i147.i1756.1 = uitofp i32 %tt2 to float --+ %div.i.i.i1749.1 = fdiv float 0.000000e+00, %conv.i147.i1756.1 --+ %cond.i.i.i1751.1 = select i1 %cmp1.i.i.i1746.1, float 0.000000e+00, float %div.i.i.i1749.1 --+ %tt3 = lshr i64 %v1, 32 --+ %2 = trunc i64 %tt3 to i32 --+ %tt1 = and i32 %2, 1 --+ %cmp1.i.i.i1746 = icmp eq i32 %tt1, 0 --+ %conv.i147.i1756 = uitofp i32 %tt1 to float --+ %div.i.i.i1749 = fdiv float 0.000000e+00, %conv.i147.i1756 --+ %cond.i.i.i1751 = select i1 %cmp1.i.i.i1746, float 0.000000e+00, float %div.i.i.i1749 --+ %3 = bitcast float %cond.i.i.i1751.3 to i32 --+ %cmp.i99.i1736.3 = icmp ult i32 %3, 1333788672 --+ %conv.i.i1743.3 = fptoui float %cond.i.i.i1751.3 to i32 --+ %4 = icmp ne i32 %conv.i.i1743.3, 0 --+ %narrow = select i1 %cmp.i99.i1736.3, i1 %4, i1 false --+ %5 = bitcast float %cond.i.i.i1751.2 to i32 --+ %cmp.i99.i1736.2 = icmp ult i32 %5, 1333788672 --+ %conv.i.i1743.2 = fptoui float %cond.i.i.i1751.2 to i32 --+ %narrow1 = select i1 %cmp.i99.i1736.2, i32 %conv.i.i1743.2, i32 0 --+ %6 = zext i1 %narrow to i32 --+ %7 = or i32 %narrow1, %6 --+ %8 = bitcast float %cond.i.i.i1751.1 to i32 --+ %cmp.i99.i1736.1 = icmp ult i32 %8, 1333788672 --+ %conv.i.i1743.1 = fptoui float %cond.i.i.i1751.1 to i32 --+ %narrow2 = select i1 %cmp.i99.i1736.1, i32 %conv.i.i1743.1, i32 0 --+ %rv3 = or i32 %7, %narrow2 --+ %9 = bitcast float %cond.i.i.i1751 to i32 --+ %cmp.i99.i1736 = icmp ult i32 %9, 1333788672 --+ %conv.i.i1743 = fptoui float %cond.i.i.i1751 to i32 --+ %narrow4 = select i1 %cmp.i99.i1736, i32 %conv.i.i1743, i32 0 --+ %rt5 = or i32 %rv3, %narrow4 --+ %rt = zext i32 %rt5 to i64 --+ store i64 %rt, ptr %v2, align 1 --+ ret i1 false --+} --diff -ruN --strip-trailing-cr a/mlir/lib/Dialect/Tensor/Utils/Utils.cpp b/mlir/lib/Dialect/Tensor/Utils/Utils.cpp ----- a/mlir/lib/Dialect/Tensor/Utils/Utils.cpp --+++ b/mlir/lib/Dialect/Tensor/Utils/Utils.cpp --@@ -16,7 +16,7 @@ -- #include "mlir/Dialect/Arith/IR/Arith.h" -- #include "mlir/Dialect/Arith/Utils/Utils.h" -- #include "mlir/Dialect/Utils/IndexingUtils.h" ---#include "mlir/Dialect/Vector/IR//VectorOps.h" --+#include "mlir/Dialect/Vector/IR/VectorOps.h" -- #include "mlir/Interfaces/ValueBoundsOpInterface.h" -- -- using namespace mlir; -+ TEST_CONSTEXPR_CXX23 bool test() { -+ // Basic test -+ { -+@@ -112,12 +117,33 @@ -+ WithNonTrivialDtor<16>, -+ WithNonTrivialDtor<256>>; -+ types::for_each(TrickyCookieTypes(), [] { -+- types::for_each(types::type_list, CustomDeleter>(), [] { -+- std::unique_ptr p(new T[3]); -++ // Array allocated with `new T[n]`, default deleter -++ { -++ std::unique_ptr> p(new T[3]); -++ assert(p[0] == T()); -++ assert(p[1] == T()); -++ assert(p[2] == T()); -++ } -++ -++ // Array allocated with `new T[n]`, custom deleter -++ { -++ std::unique_ptr> p(new T[3]); -++ assert(p[0] == T()); -++ assert(p[1] == T()); -++ assert(p[2] == T()); -++ } -++ -++ // Array not allocated with `new T[n]`, custom deleter -++ // -++ // This test aims to ensure that the implementation doesn't try to use an array cookie -++ // when there is none. -++ { -++ T array[50] = {}; -++ std::unique_ptr p(&array[0]); -+ assert(p[0] == T()); -+ assert(p[1] == T()); -+ assert(p[2] == T()); -+- }); -++ } -+ }); ++ // These variables are initialized by Writer and should not be used before ++ // Writer is initialized. ++diff -ruN --strip-trailing-cr a/lld/ELF/Driver.cpp b/lld/ELF/Driver.cpp ++--- a/lld/ELF/Driver.cpp +++++ b/lld/ELF/Driver.cpp ++@@ -99,7 +99,7 @@ ++ driver.~LinkerDriver(); ++ new (&driver) LinkerDriver(*this); ++ script = nullptr; ++- target.reset(); +++ target = nullptr; + +- template class basic_string { +- public: +-@@ -89,12 +77,7 @@ +- }; +- #else // !SUBCLASS_PADDING ++ bufferStart = nullptr; ++ mainPart = nullptr; ++@@ -3126,7 +3126,7 @@ ++ // The Target instance handles target-specific stuff, such as applying ++ // relocations or writing a PLT section. It also contains target-dependent ++ // values such as a default image base address. ++- setTarget(ctx); +++ ctx.target = getTarget(ctx); + +--#ifdef NON_STANDARD_PADDING +- unsigned char __padding[sizeof(value_type) - 1]; +--#else +-- [[no_unique_address]] __padding __padding_; +--#endif +-- +- #ifdef BITMASKS +- unsigned char __size_; +- #else // !BITMASKS +-@@ -146,26 +129,21 @@ +- union { +- #ifdef BITMASKS +- unsigned char __size_; +--#else // !BITMASKS +-+#else +- struct { +- unsigned char __is_long_ : 1; +- unsigned char __size_ : 7; +- }; +--#endif // BITMASKS +-+#endif +- value_type __lx; +- }; +--#else // !SHORT_UNION +-+#else +- BEGIN_PACKED_ANON_STRUCT +- unsigned char __is_long_ : 1; +- unsigned char __size_ : 7; +- END_PACKED_ANON_STRUCT +--#ifdef NON_STANDARD_PADDING +-- unsigned char __padding[sizeof(value_type) - 1]; +--#else // !NON_STANDARD_PADDING +-- [[no_unique_address]] __padding __padding_; +--#endif // NON_STANDARD_PADDING +-- +--#endif // SHORT_UNION +-+ char __padding_[sizeof(value_type) - 1]; +-+#endif +- value_type __data_[__min_cap]; +- }; ++ ctx.arg.eflags = ctx.target->calcEFlags(); ++ // maxPageSize (sometimes called abi page size) is the maximum page size that ++diff -ruN --strip-trailing-cr a/lld/ELF/Target.cpp b/lld/ELF/Target.cpp ++--- a/lld/ELF/Target.cpp +++++ b/lld/ELF/Target.cpp ++@@ -45,39 +45,39 @@ ++ return std::string(s); ++ } + +-diff -ruN --strip-trailing-cr a/lldb/test/API/functionalities/data-formatter/data-formatter-stl/libcxx-simulators/string/TestDataFormatterLibcxxStringSimulator.py b/lldb/test/API/functionalities/data-formatter/data-formatter-stl/libcxx-simulators/string/TestDataFormatterLibcxxStringSimulator.py +---- a/lldb/test/API/functionalities/data-formatter/data-formatter-stl/libcxx-simulators/string/TestDataFormatterLibcxxStringSimulator.py +-+++ b/lldb/test/API/functionalities/data-formatter/data-formatter-stl/libcxx-simulators/string/TestDataFormatterLibcxxStringSimulator.py +-@@ -27,7 +27,7 @@ ++-void elf::setTarget(Ctx &ctx) { +++TargetInfo *elf::getTarget(Ctx &ctx) { ++ switch (ctx.arg.emachine) { ++ case EM_386: ++ case EM_IAMCU: ++- return setX86TargetInfo(ctx); +++ return getX86TargetInfo(ctx); ++ case EM_AARCH64: ++- return setAArch64TargetInfo(ctx); +++ return getAArch64TargetInfo(ctx); ++ case EM_AMDGPU: ++- return setAMDGPUTargetInfo(ctx); +++ return getAMDGPUTargetInfo(ctx); ++ case EM_ARM: ++- return setARMTargetInfo(ctx); +++ return getARMTargetInfo(ctx); ++ case EM_AVR: ++- return setAVRTargetInfo(ctx); +++ return getAVRTargetInfo(ctx); ++ case EM_HEXAGON: ++- return setHexagonTargetInfo(ctx); +++ return getHexagonTargetInfo(ctx); ++ case EM_LOONGARCH: ++- return setLoongArchTargetInfo(ctx); +++ return getLoongArchTargetInfo(ctx); ++ case EM_MIPS: ++- return setMipsTargetInfo(ctx); +++ return getMipsTargetInfo(ctx); ++ case EM_MSP430: ++- return setMSP430TargetInfo(ctx); +++ return getMSP430TargetInfo(ctx); ++ case EM_PPC: ++- return setPPCTargetInfo(ctx); +++ return getPPCTargetInfo(ctx); ++ case EM_PPC64: ++- return setPPC64TargetInfo(ctx); +++ return getPPC64TargetInfo(ctx); ++ case EM_RISCV: ++- return setRISCVTargetInfo(ctx); +++ return getRISCVTargetInfo(ctx); ++ case EM_SPARCV9: ++- return setSPARCV9TargetInfo(ctx); +++ return getSPARCV9TargetInfo(ctx); ++ case EM_S390: ++- return setSystemZTargetInfo(ctx); +++ return getSystemZTargetInfo(ctx); ++ case EM_X86_64: ++- return setX86_64TargetInfo(ctx); +++ return getX86_64TargetInfo(ctx); ++ default: ++ fatal("unsupported e_machine value: " + Twine(ctx.arg.emachine)); + } -+ #endif // C++20 -+diff -ruN --strip-trailing-cr a/lldb/test/API/functionalities/data-formatter/data-formatter-stl/libcxx-simulators/string/main.cpp b/lldb/test/API/functionalities/data-formatter/data-formatter-stl/libcxx-simulators/string/main.cpp -+--- a/lldb/test/API/functionalities/data-formatter/data-formatter-stl/libcxx-simulators/string/main.cpp -++++ b/lldb/test/API/functionalities/data-formatter/data-formatter-stl/libcxx-simulators/string/main.cpp -+@@ -20,11 +20,7 @@ -+ // Pre-D128285 layout. -+ #define PACKED_ANON_STRUCT -+ #endif -+-#if REVISION <= 4 -+-// Pre-2a1ef74 layout. -+-#define NON_STANDARD_PADDING -+-#endif -+-// REVISION == 5: current layout -++// REVISION == 4: current layout -+ -+ #ifdef PACKED_ANON_STRUCT -+ #define BEGIN_PACKED_ANON_STRUCT struct __attribute__((packed)) { -+@@ -38,7 +34,6 @@ -+ namespace std { -+ namespace __lldb { -+ -+-#ifdef NON_STANDARD_PADDING -+ #if defined(ALTERNATE_LAYOUT) && defined(SUBCLASS_PADDING) -+ template struct __padding { -+ unsigned char __xx[sizeof(_CharT) - 1]; -+@@ -46,13 +41,6 @@ -+ -+ template struct __padding<_CharT, 1> {}; -+ #endif -+-#else // !NON_STANDARD_PADDING -+-template struct __padding { -+- char __padding_[_PaddingSize]; -+-}; -+- -+-template <> struct __padding<0> {}; -+-#endif -+ -+ template class basic_string { -+ public: -+@@ -89,12 +77,7 @@ -+ }; -+ #else // !SUBCLASS_PADDING -+ -+-#ifdef NON_STANDARD_PADDING -+ unsigned char __padding[sizeof(value_type) - 1]; -+-#else -+- [[no_unique_address]] __padding __padding_; -+-#endif -+- -+ #ifdef BITMASKS -+ unsigned char __size_; -+ #else // !BITMASKS -+@@ -146,26 +129,21 @@ -+ union { -+ #ifdef BITMASKS -+ unsigned char __size_; -+-#else // !BITMASKS -++#else -+ struct { -+ unsigned char __is_long_ : 1; -+ unsigned char __size_ : 7; -+ }; -+-#endif // BITMASKS -++#endif -+ value_type __lx; -+ }; -+-#else // !SHORT_UNION -++#else -+ BEGIN_PACKED_ANON_STRUCT -+ unsigned char __is_long_ : 1; -+ unsigned char __size_ : 7; -+ END_PACKED_ANON_STRUCT -+-#ifdef NON_STANDARD_PADDING -+- unsigned char __padding[sizeof(value_type) - 1]; -+-#else // !NON_STANDARD_PADDING -+- [[no_unique_address]] __padding __padding_; -+-#endif // NON_STANDARD_PADDING -+- -+-#endif // SHORT_UNION -++ char __padding_[sizeof(value_type) - 1]; -++#endif -+ value_type __data_[__min_cap]; -+ }; -+ -+diff -ruN --strip-trailing-cr a/lldb/test/API/functionalities/data-formatter/data-formatter-stl/libcxx-simulators/string/TestDataFormatterLibcxxStringSimulator.py b/lldb/test/API/functionalities/data-formatter/data-formatter-stl/libcxx-simulators/string/TestDataFormatterLibcxxStringSimulator.py -+--- a/lldb/test/API/functionalities/data-formatter/data-formatter-stl/libcxx-simulators/string/TestDataFormatterLibcxxStringSimulator.py -++++ b/lldb/test/API/functionalities/data-formatter/data-formatter-stl/libcxx-simulators/string/TestDataFormatterLibcxxStringSimulator.py -+@@ -27,7 +27,7 @@ -+ -+ -+ for v in [None, "ALTERNATE_LAYOUT"]: -+- for r in range(6): -++ for r in range(5): -+ for c in range(3): -+ name = "test_r%d_c%d" % (r, c) -+ defines = ["REVISION=%d" % r, "COMPRESSED_PAIR_REV=%d" % c] - diff -ruN --strip-trailing-cr a/utils/bazel/llvm-project-overlay/libc/BUILD.bazel b/utils/bazel/llvm-project-overlay/libc/BUILD.bazel - --- a/utils/bazel/llvm-project-overlay/libc/BUILD.bazel - +++ b/utils/bazel/llvm-project-overlay/libc/BUILD.bazel -@@ -175,33 +358,3 @@ diff -ruN --strip-trailing-cr a/utils/bazel/llvm-project-overlay/libc/BUILD.baze - ) ++diff -ruN --strip-trailing-cr a/lld/ELF/Target.h b/lld/ELF/Target.h ++--- a/lld/ELF/Target.h +++++ b/lld/ELF/Target.h ++@@ -179,21 +179,21 @@ ++ uint64_t defaultImageBase = 0x10000; ++ }; + ++-void setAArch64TargetInfo(Ctx &); ++-void setAMDGPUTargetInfo(Ctx &); ++-void setARMTargetInfo(Ctx &); ++-void setAVRTargetInfo(Ctx &); ++-void setHexagonTargetInfo(Ctx &); ++-void setLoongArchTargetInfo(Ctx &); ++-void setMSP430TargetInfo(Ctx &); ++-void setMipsTargetInfo(Ctx &); ++-void setPPC64TargetInfo(Ctx &); ++-void setPPCTargetInfo(Ctx &); ++-void setRISCVTargetInfo(Ctx &); ++-void setSPARCV9TargetInfo(Ctx &); ++-void setSystemZTargetInfo(Ctx &); ++-void setX86TargetInfo(Ctx &); ++-void setX86_64TargetInfo(Ctx &); +++TargetInfo *getAArch64TargetInfo(Ctx &); +++TargetInfo *getAMDGPUTargetInfo(Ctx &); +++TargetInfo *getARMTargetInfo(Ctx &); +++TargetInfo *getAVRTargetInfo(Ctx &); +++TargetInfo *getHexagonTargetInfo(Ctx &); +++TargetInfo *getLoongArchTargetInfo(Ctx &); +++TargetInfo *getMSP430TargetInfo(Ctx &); +++TargetInfo *getMipsTargetInfo(Ctx &); +++TargetInfo *getPPC64TargetInfo(Ctx &); +++TargetInfo *getPPCTargetInfo(Ctx &); +++TargetInfo *getRISCVTargetInfo(Ctx &); +++TargetInfo *getSPARCV9TargetInfo(Ctx &); +++TargetInfo *getSystemZTargetInfo(Ctx &); +++TargetInfo *getX86TargetInfo(Ctx &); +++TargetInfo *getX86_64TargetInfo(Ctx &); - libc_support_library( --diff -ruN --strip-trailing-cr a/utils/bazel/llvm-project-overlay/lldb/source/Plugins/BUILD.bazel b/utils/bazel/llvm-project-overlay/lldb/source/Plugins/BUILD.bazel ----- a/utils/bazel/llvm-project-overlay/lldb/source/Plugins/BUILD.bazel --+++ b/utils/bazel/llvm-project-overlay/lldb/source/Plugins/BUILD.bazel --@@ -1091,6 +1091,7 @@ -- "//lldb:Utility", -- "//lldb:UtilityPrivateHeaders", -- "//llvm:Core", --+ "//llvm:Support", -- "//llvm:TargetParser", -- ], +- for v in [None, "ALTERNATE_LAYOUT"]: +-- for r in range(6): +-+ for r in range(5): +- for c in range(3): +- name = "test_r%d_c%d" % (r, c) +- defines = ["REVISION=%d" % r, "COMPRESSED_PAIR_REV=%d" % c] +-diff -ruN --strip-trailing-cr a/utils/bazel/llvm-project-overlay/libc/BUILD.bazel b/utils/bazel/llvm-project-overlay/libc/BUILD.bazel +---- a/utils/bazel/llvm-project-overlay/libc/BUILD.bazel +-+++ b/utils/bazel/llvm-project-overlay/libc/BUILD.bazel +-@@ -148,6 +148,9 @@ +- libc_support_library( +- name = "hdr_stdio_macros", +- hdrs = ["hdr/stdio_macros.h"], +-+ deps = [ +-+ ":hdr_stdio_overlay", +-+ ], - ) --diff -ruN --strip-trailing-cr a/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel b/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel ----- a/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel --+++ b/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel --@@ -7694,6 +7694,7 @@ -- ":DialectUtils", -- ":TensorDialect", -- ":ValueBoundsOpInterface", --+ ":VectorDialect", -- ], ++ struct ErrorPlace { ++ InputSectionBase *isec; ++@@ -251,7 +251,7 @@ ++ void createTaggedSymbols(Ctx &); ++ void initSymbolAnchors(Ctx &); + +- libc_support_library( +-@@ -215,11 +218,17 @@ +- libc_support_library( +- name = "types_off_t", +- hdrs = ["hdr/types/off_t.h"], +-+ deps = [ +-+ ":hdr_stdio_overlay", +-+ ], - ) -- --@@ -13902,6 +13903,7 @@ -- ":BufferViewFlowOpInterfaceIncGen", -- ":BufferizableOpInterfaceIncGen", -- ":BufferizationEnumsIncGen", --+ ":FunctionInterfaces", -- ":IR", -- ":Support", -- "//llvm:Support", ++-void setTarget(Ctx &); +++TargetInfo *getTarget(Ctx &); + +- libc_support_library( +- name = "types_FILE", +- hdrs = ["hdr/types/FILE.h"], +-+ deps = [ +-+ ":hdr_stdio_overlay", +-+ ], +- ) ++ template bool isMipsPIC(const Defined *sym); + +- libc_support_library( diff --git a/third_party/llvm/workspace.bzl b/third_party/llvm/workspace.bzl -index 6b843a7..50e3160 100644 +index 50e3160..e939a2e 100644 --- a/third_party/llvm/workspace.bzl +++ b/third_party/llvm/workspace.bzl @@ -4,8 +4,8 @@ load("//third_party:repo.bzl", "tf_http_archive") def repo(name): """Imports LLVM.""" -- LLVM_COMMIT = "68c04b0ae62d8431d72d8b47fc13008002ee4387" -- LLVM_SHA256 = "90b9b631d6670a99d22e022d180f2877bf523cf9c6dc05876a6912dfd1a3e332" -+ LLVM_COMMIT = "82f5acfbec65e1a645d902f746253eeaf0bd2d70" -+ LLVM_SHA256 = "c621acf6f202cfcbd0c73f381272b70c51c95f0ecb62c17d2d5fe0156a80082a" +- LLVM_COMMIT = "82f5acfbec65e1a645d902f746253eeaf0bd2d70" +- LLVM_SHA256 = "c621acf6f202cfcbd0c73f381272b70c51c95f0ecb62c17d2d5fe0156a80082a" ++ LLVM_COMMIT = "2918e779a9545a66c0031b03b3af5bf4d8517cec" ++ LLVM_SHA256 = "efd6ac33aba00a6eac0047533a99d6d766dc9dae9ec5ba8e74ff297bfd4c9f0e" tf_http_archive( name = name, diff --git a/third_party/shardy/workspace.bzl b/third_party/shardy/workspace.bzl index d4edca2d3a645..25454dbc21375 100644 --- a/third_party/shardy/workspace.bzl +++ b/third_party/shardy/workspace.bzl @@ -3,8 +3,8 @@ load("//third_party:repo.bzl", "tf_http_archive", "tf_mirror_urls") def repo(): - SHARDY_COMMIT = "efbd41448d245c78c6cbd73689d1a6fb40839152" - SHARDY_SHA256 = "c746965803e9bfb4ccf20366eae60fbe32fcc5797520dfe0657527d5f47f6768" + SHARDY_COMMIT = "a21e233b5c6cc45c806d3e9e7960d8b1f0fbab4d" + SHARDY_SHA256 = "db7aa65c3386c2560f43faf1eae2e7b824282624defcf76f4901ec8bda134c1f" tf_http_archive( name = "shardy", diff --git a/third_party/triton/llvm_integration/cl683501567.patch b/third_party/triton/llvm_integration/cl683501567.patch new file mode 100644 index 0000000000000..7395934253fc0 --- /dev/null +++ b/third_party/triton/llvm_integration/cl683501567.patch @@ -0,0 +1,13 @@ + +--- a/lib/Target/LLVMIR/LLVMDIScope.cpp 2024-09-16 13:44:40.000000000 -0700 ++++ b/lib/Target/LLVMIR/LLVMDIScope.cpp 2024-10-08 22:38:50.000000000 -0700 +@@ -104,7 +104,8 @@ + auto subprogramAttr = LLVM::DISubprogramAttr::get( + context, distinctId, compileUnitAttr, fileAttr, funcNameAttr, + funcNameAttr, fileAttr, /*line=*/line, /*scopeline=*/line, +- subprogramFlags, subroutineTypeAttr, /*retainNodes=*/{}); ++ subprogramFlags, subroutineTypeAttr, /*retainNodes=*/{}, ++ /*annotations=*/{}); + funcOp->setLoc(FusedLoc::get(context, {loc}, subprogramAttr)); + } + diff --git a/third_party/triton/llvm_integration/series.bzl b/third_party/triton/llvm_integration/series.bzl index ae99559a93cc4..2ddfcdd22989b 100644 --- a/third_party/triton/llvm_integration/series.bzl +++ b/third_party/triton/llvm_integration/series.bzl @@ -9,5 +9,6 @@ LLVM nor MLIR integrator, please do not add any patches to this list. llvm_patch_list = [ "//third_party/triton:llvm_integration/cl680875920.patch", + "//third_party/triton:llvm_integration/cl683501567.patch", # Add new patches just above this line ] diff --git a/third_party/tsl/third_party/llvm/generated.patch b/third_party/tsl/third_party/llvm/generated.patch index f8fa2d147648e..943d396a43efd 100644 --- a/third_party/tsl/third_party/llvm/generated.patch +++ b/third_party/tsl/third_party/llvm/generated.patch @@ -1,360 +1,419 @@ Auto generated patch. Do not edit or delete it, even if empty. -diff -ruN --strip-trailing-cr a/libcxx/include/__memory/array_cookie.h b/libcxx/include/__memory/array_cookie.h ---- a/libcxx/include/__memory/array_cookie.h -+++ b/libcxx/include/__memory/array_cookie.h -@@ -24,7 +24,7 @@ - _LIBCPP_BEGIN_NAMESPACE_STD +diff -ruN --strip-trailing-cr a/clang/test/SemaCXX/pragma-pack-packed-2.cpp b/clang/test/SemaCXX/pragma-pack-packed-2.cpp +--- a/clang/test/SemaCXX/pragma-pack-packed-2.cpp ++++ b/clang/test/SemaCXX/pragma-pack-packed-2.cpp +@@ -1,6 +1,6 @@ + // RUN: %clang_cc1 -triple s390x-ibm-zos -fzos-extensions -fsyntax-only -verify %s + // RUN: %clang_cc1 -triple s390x-ibm-zos -fzos-extensions -fxl-pragma-pack -fsyntax-only -verify %s +-// RUN: %clang -target s390x-ibm-zos -S -emit-llvm -Xclang -verify -fno-xl-pragma-pack %s ++// RUN: %clang -target s390x-ibm-zos -S -emit-llvm -Xclang -verify -fno-xl-pragma-pack -o %t.ll %s - // Trait representing whether a type requires an array cookie at the start of its allocation when --// allocated as `new T[n]` and deallocated as `delete array`. -+// allocated as `new T[n]` and deallocated as `delete[] array`. - // - // Under the Itanium C++ ABI [1], we know that an array cookie is available unless `T` is trivially - // destructible and the call to `operator delete[]` is not a sized operator delete. Under ABIs other -diff -ruN --strip-trailing-cr a/libcxx/include/__memory/unique_ptr.h b/libcxx/include/__memory/unique_ptr.h ---- a/libcxx/include/__memory/unique_ptr.h -+++ b/libcxx/include/__memory/unique_ptr.h -@@ -102,6 +102,12 @@ - }; + #pragma pack(show) // expected-warning {{value of #pragma pack(show) == 8}} + #pragma pack(twobyte) +diff -ruN --strip-trailing-cr a/lld/ELF/Arch/AArch64.cpp b/lld/ELF/Arch/AArch64.cpp +--- a/lld/ELF/Arch/AArch64.cpp ++++ b/lld/ELF/Arch/AArch64.cpp +@@ -1208,10 +1208,12 @@ + } + } - template -+struct __is_default_deleter : false_type {}; -+ -+template -+struct __is_default_deleter > : true_type {}; -+ -+template - struct __unique_ptr_deleter_sfinae { - static_assert(!is_reference<_Deleter>::value, "incorrect specialization"); - typedef const _Deleter& __lval_ref_type; -@@ -307,11 +313,16 @@ - // 1. When an array cookie (see [1]) exists at the beginning of the array allocation, we are - // able to reuse that cookie to extract the size of the array and perform bounds checking. - // An array cookie is a size inserted at the beginning of the allocation by the compiler. --// That size is inserted implicitly when doing `new T[n]` in some cases, and its purpose --// is to allow the runtime to destroy the `n` array elements when doing `delete array`. -+// That size is inserted implicitly when doing `new T[n]` in some cases (as of writing this -+// exactly when the array elements are not trivially destructible), and its main purpose is -+// to allow the runtime to destroy the `n` array elements when doing `delete[] array`. - // When we are able to use array cookies, we reuse information already available in the - // current runtime, so bounds checking does not require changing libc++'s ABI. - // -+// However, note that we cannot assume the presence of an array cookie when a custom deleter -+// is used, because the unique_ptr could have been created from an allocation that wasn't -+// obtained via `new T[n]` (since it may not be deleted with `delete[] arr`). -+// - // 2. When the "bounded unique_ptr" ABI configuration (controlled by `_LIBCPP_ABI_BOUNDED_UNIQUE_PTR`) - // is enabled, we store the size of the allocation (when it is known) so we can check it when - // indexing into the `unique_ptr`. That changes the layout of `std::unique_ptr`, which is -@@ -328,7 +339,7 @@ - // try to fall back to using an array cookie when available. - // - // Finally, note that when this ABI configuration is enabled, we have no choice but to always --// make space for a size to be stored in the unique_ptr. Indeed, while we might want to avoid -+// make space for the size to be stored in the unique_ptr. Indeed, while we might want to avoid - // storing the size when an array cookie is available, knowing whether an array cookie is available - // requires the type stored in the unique_ptr to be complete, while unique_ptr can normally - // accommodate incomplete types. -@@ -339,7 +350,9 @@ - __unique_ptr_array_bounds_stateless() = default; - _LIBCPP_HIDE_FROM_ABI _LIBCPP_CONSTEXPR explicit __unique_ptr_array_bounds_stateless(size_t) {} +-void elf::setAArch64TargetInfo(Ctx &ctx) { ++TargetInfo *elf::getAArch64TargetInfo(Ctx &ctx) { + if ((ctx.arg.andFeatures & GNU_PROPERTY_AARCH64_FEATURE_1_BTI) || +- ctx.arg.zPacPlt) +- ctx.target.reset(new AArch64BtiPac(ctx)); +- else +- ctx.target.reset(new AArch64(ctx)); ++ ctx.arg.zPacPlt) { ++ static AArch64BtiPac t(ctx); ++ return &t; ++ } ++ static AArch64 t(ctx); ++ return &t; + } +diff -ruN --strip-trailing-cr a/lld/ELF/Arch/AMDGPU.cpp b/lld/ELF/Arch/AMDGPU.cpp +--- a/lld/ELF/Arch/AMDGPU.cpp ++++ b/lld/ELF/Arch/AMDGPU.cpp +@@ -219,4 +219,7 @@ + } + } + +-void elf::setAMDGPUTargetInfo(Ctx &ctx) { ctx.target.reset(new AMDGPU(ctx)); } ++TargetInfo *elf::getAMDGPUTargetInfo(Ctx &ctx) { ++ static AMDGPU target(ctx); ++ return ⌖ ++} +diff -ruN --strip-trailing-cr a/lld/ELF/Arch/ARM.cpp b/lld/ELF/Arch/ARM.cpp +--- a/lld/ELF/Arch/ARM.cpp ++++ b/lld/ELF/Arch/ARM.cpp +@@ -1533,7 +1533,10 @@ + "': " + toString(std::move(e))); + } + +-void elf::setARMTargetInfo(Ctx &ctx) { ctx.target.reset(new ARM(ctx)); } ++TargetInfo *elf::getARMTargetInfo(Ctx &ctx) { ++ static ARM target(ctx); ++ return ⌖ ++} -- template ::value, int> = 0> -+ template ::value && __has_array_cookie<_Tp>::value, int> = 0> - _LIBCPP_HIDE_FROM_ABI _LIBCPP_CONSTEXPR bool __in_bounds(_Tp* __ptr, size_t __index) const { - // In constant expressions, we can't check the array cookie so we just pretend that the index - // is in-bounds. The compiler catches invalid accesses anyway. -@@ -349,7 +362,9 @@ - return __index < __cookie; + template void elf::writeARMCmseImportLib(Ctx &); + template void elf::writeARMCmseImportLib(Ctx &); +diff -ruN --strip-trailing-cr a/lld/ELF/Arch/AVR.cpp b/lld/ELF/Arch/AVR.cpp +--- a/lld/ELF/Arch/AVR.cpp ++++ b/lld/ELF/Arch/AVR.cpp +@@ -267,7 +267,10 @@ } + } -- template ::value, int> = 0> -+ template ::value || !__has_array_cookie<_Tp>::value, int> = 0> - _LIBCPP_HIDE_FROM_ABI _LIBCPP_CONSTEXPR bool __in_bounds(_Tp*, size_t) const { - return true; // If we don't have an array cookie, we assume the access is in-bounds +-void elf::setAVRTargetInfo(Ctx &ctx) { ctx.target.reset(new AVR(ctx)); } ++TargetInfo *elf::getAVRTargetInfo(Ctx &ctx) { ++ static AVR target(ctx); ++ return ⌖ ++} + + static uint32_t getEFlags(InputFile *file) { + return cast>(file)->getObj().getHeader().e_flags; +diff -ruN --strip-trailing-cr a/lld/ELF/Arch/Hexagon.cpp b/lld/ELF/Arch/Hexagon.cpp +--- a/lld/ELF/Arch/Hexagon.cpp ++++ b/lld/ELF/Arch/Hexagon.cpp +@@ -404,4 +404,7 @@ } -@@ -365,7 +380,9 @@ - _LIBCPP_HIDE_FROM_ABI _LIBCPP_CONSTEXPR explicit __unique_ptr_array_bounds_stored(size_t __size) : __size_(__size) {} + } - // Use the array cookie if there's one -- template ::value, int> = 0> -+ template ::value && __has_array_cookie<_Tp>::value, int> = 0> - _LIBCPP_HIDE_FROM_ABI _LIBCPP_CONSTEXPR bool __in_bounds(_Tp* __ptr, size_t __index) const { - if (__libcpp_is_constant_evaluated()) - return true; -@@ -374,7 +391,9 @@ +-void elf::setHexagonTargetInfo(Ctx &ctx) { ctx.target.reset(new Hexagon(ctx)); } ++TargetInfo *elf::getHexagonTargetInfo(Ctx &ctx) { ++ static Hexagon target(ctx); ++ return ⌖ ++} +diff -ruN --strip-trailing-cr a/lld/ELF/Arch/LoongArch.cpp b/lld/ELF/Arch/LoongArch.cpp +--- a/lld/ELF/Arch/LoongArch.cpp ++++ b/lld/ELF/Arch/LoongArch.cpp +@@ -893,6 +893,7 @@ } + } + +-void elf::setLoongArchTargetInfo(Ctx &ctx) { +- ctx.target.reset(new LoongArch(ctx)); ++TargetInfo *elf::getLoongArchTargetInfo(Ctx &ctx) { ++ static LoongArch target(ctx); ++ return ⌖ + } +diff -ruN --strip-trailing-cr a/lld/ELF/Arch/Mips.cpp b/lld/ELF/Arch/Mips.cpp +--- a/lld/ELF/Arch/Mips.cpp ++++ b/lld/ELF/Arch/Mips.cpp +@@ -779,23 +779,23 @@ + return cast>(file)->getObj().getHeader().e_flags & EF_MIPS_PIC; + } - // Otherwise, fall back on the stored size (if any) -- template ::value, int> = 0> -+ template ::value || !__has_array_cookie<_Tp>::value, int> = 0> - _LIBCPP_HIDE_FROM_ABI _LIBCPP_CONSTEXPR bool __in_bounds(_Tp*, size_t __index) const { - return __index < __size_; +-void elf::setMipsTargetInfo(Ctx &ctx) { ++TargetInfo *elf::getMipsTargetInfo(Ctx &ctx) { + switch (ctx.arg.ekind) { + case ELF32LEKind: { +- ctx.target.reset(new MIPS(ctx)); +- return; ++ static MIPS t(ctx); ++ return &t; } -@@ -562,7 +581,7 @@ + case ELF32BEKind: { +- ctx.target.reset(new MIPS(ctx)); +- return; ++ static MIPS t(ctx); ++ return &t; } - - _LIBCPP_HIDE_FROM_ABI _LIBCPP_CONSTEXPR_SINCE_CXX23 __add_lvalue_reference_t<_Tp> operator[](size_t __i) const { -- _LIBCPP_ASSERT_VALID_ELEMENT_ACCESS(__checker_.__in_bounds(std::__to_address(__ptr_), __i), -+ _LIBCPP_ASSERT_VALID_ELEMENT_ACCESS(__checker_.__in_bounds(std::__to_address(__ptr_), __i), - "unique_ptr::operator[](index): index out of range"); - return __ptr_[__i]; + case ELF64LEKind: { +- ctx.target.reset(new MIPS(ctx)); +- return; ++ static MIPS t(ctx); ++ return &t; + } + case ELF64BEKind: { +- ctx.target.reset(new MIPS(ctx)); +- return; ++ static MIPS t(ctx); ++ return &t; + } + default: + llvm_unreachable("unsupported target"); +diff -ruN --strip-trailing-cr a/lld/ELF/Arch/MSP430.cpp b/lld/ELF/Arch/MSP430.cpp +--- a/lld/ELF/Arch/MSP430.cpp ++++ b/lld/ELF/Arch/MSP430.cpp +@@ -88,4 +88,7 @@ } -diff -ruN --strip-trailing-cr a/libcxx/test/std/utilities/smartptr/unique.ptr/unique.ptr.class/unique.ptr.observers/assert.subscript.pass.cpp b/libcxx/test/std/utilities/smartptr/unique.ptr/unique.ptr.class/unique.ptr.observers/assert.subscript.pass.cpp ---- a/libcxx/test/std/utilities/smartptr/unique.ptr/unique.ptr.class/unique.ptr.observers/assert.subscript.pass.cpp -+++ b/libcxx/test/std/utilities/smartptr/unique.ptr/unique.ptr.class/unique.ptr.observers/assert.subscript.pass.cpp -@@ -48,31 +48,24 @@ + } - template - void test() { -- // For types with an array cookie, we can always detect OOB accesses. -+ // For types with an array cookie, we can always detect OOB accesses. Note that reliance on an array -+ // cookie is limited to the default deleter, since a unique_ptr with a custom deleter may not have -+ // been allocated with `new T[n]`. - { -- // Check with the default deleter - { -- { -- std::unique_ptr ptr(new WithCookie[5]); -- TEST_LIBCPP_ASSERT_FAILURE(ptr[6], "unique_ptr::operator[](index): index out of range"); -- } -- { -- std::unique_ptr ptr = std::make_unique(5); -- TEST_LIBCPP_ASSERT_FAILURE(ptr[6], "unique_ptr::operator[](index): index out of range"); -- } --#if TEST_STD_VER >= 20 -- { -- std::unique_ptr ptr = std::make_unique_for_overwrite(5); -- TEST_LIBCPP_ASSERT_FAILURE(ptr[6] = WithCookie(), "unique_ptr::operator[](index): index out of range"); -- } --#endif -+ std::unique_ptr ptr(new WithCookie[5]); -+ TEST_LIBCPP_ASSERT_FAILURE(ptr[6], "unique_ptr::operator[](index): index out of range"); - } -- -- // Check with a custom deleter - { -- std::unique_ptr ptr(new WithCookie[5]); -+ std::unique_ptr ptr = std::make_unique(5); - TEST_LIBCPP_ASSERT_FAILURE(ptr[6], "unique_ptr::operator[](index): index out of range"); - } -+#if TEST_STD_VER >= 20 -+ { -+ std::unique_ptr ptr = std::make_unique_for_overwrite(5); -+ TEST_LIBCPP_ASSERT_FAILURE(ptr[6] = WithCookie(), "unique_ptr::operator[](index): index out of range"); -+ } -+#endif +-void elf::setMSP430TargetInfo(Ctx &ctx) { ctx.target.reset(new MSP430(ctx)); } ++TargetInfo *elf::getMSP430TargetInfo(Ctx &ctx) { ++ static MSP430 target(ctx); ++ return ⌖ ++} +diff -ruN --strip-trailing-cr a/lld/ELF/Arch/PPC64.cpp b/lld/ELF/Arch/PPC64.cpp +--- a/lld/ELF/Arch/PPC64.cpp ++++ b/lld/ELF/Arch/PPC64.cpp +@@ -1747,4 +1747,7 @@ + return true; + } + +-void elf::setPPC64TargetInfo(Ctx &ctx) { ctx.target.reset(new PPC64(ctx)); } ++TargetInfo *elf::getPPC64TargetInfo(Ctx &ctx) { ++ static PPC64 target(ctx); ++ return ⌖ ++} +diff -ruN --strip-trailing-cr a/lld/ELF/Arch/PPC.cpp b/lld/ELF/Arch/PPC.cpp +--- a/lld/ELF/Arch/PPC.cpp ++++ b/lld/ELF/Arch/PPC.cpp +@@ -523,4 +523,7 @@ } + } - // For types that don't have an array cookie, things are a bit more complicated. We can detect OOB accesses -@@ -97,14 +90,9 @@ - #endif +-void elf::setPPCTargetInfo(Ctx &ctx) { ctx.target.reset(new PPC(ctx)); } ++TargetInfo *elf::getPPCTargetInfo(Ctx &ctx) { ++ static PPC target(ctx); ++ return ⌖ ++} +diff -ruN --strip-trailing-cr a/lld/ELF/Arch/RISCV.cpp b/lld/ELF/Arch/RISCV.cpp +--- a/lld/ELF/Arch/RISCV.cpp ++++ b/lld/ELF/Arch/RISCV.cpp +@@ -1329,4 +1329,7 @@ + mergeAttributesSection(ctx, sections)); + } + +-void elf::setRISCVTargetInfo(Ctx &ctx) { ctx.target.reset(new RISCV(ctx)); } ++TargetInfo *elf::getRISCVTargetInfo(Ctx &ctx) { ++ static RISCV target(ctx); ++ return ⌖ ++} +diff -ruN --strip-trailing-cr a/lld/ELF/Arch/SPARCV9.cpp b/lld/ELF/Arch/SPARCV9.cpp +--- a/lld/ELF/Arch/SPARCV9.cpp ++++ b/lld/ELF/Arch/SPARCV9.cpp +@@ -193,4 +193,7 @@ + relocateNoSym(buf + 4, R_SPARC_WDISP19, -(off + 4 - pltEntrySize)); + } - // Make sure that we carry the bounds information properly through conversions, assignments, etc. -- // These tests are mostly relevant when the ABI setting is enabled (with a stateful bounds-checker), -- // but we still run them for types with an array cookie either way. -+ // These tests are only relevant when the ABI setting is enabled (with a stateful bounds-checker). - #if defined(_LIBCPP_ABI_BOUNDED_UNIQUE_PTR) -- using Types = types::type_list; --#else -- using Types = types::type_list; --#endif -- types::for_each(Types(), [] { -+ types::for_each(types::type_list(), [] { - // Bounds carried through move construction - { - std::unique_ptr ptr = std::make_unique(5); -@@ -135,6 +123,7 @@ - TEST_LIBCPP_ASSERT_FAILURE(other[6], "unique_ptr::operator[](index): index out of range"); - } - }); -+#endif +-void elf::setSPARCV9TargetInfo(Ctx &ctx) { ctx.target.reset(new SPARCV9(ctx)); } ++TargetInfo *elf::getSPARCV9TargetInfo(Ctx &ctx) { ++ static SPARCV9 target(ctx); ++ return ⌖ ++} +diff -ruN --strip-trailing-cr a/lld/ELF/Arch/SystemZ.cpp b/lld/ELF/Arch/SystemZ.cpp +--- a/lld/ELF/Arch/SystemZ.cpp ++++ b/lld/ELF/Arch/SystemZ.cpp +@@ -600,4 +600,7 @@ + } } - template -diff -ruN --strip-trailing-cr a/libcxx/test/std/utilities/smartptr/unique.ptr/unique.ptr.class/unique.ptr.observers/op_subscript.runtime.pass.cpp b/libcxx/test/std/utilities/smartptr/unique.ptr/unique.ptr.class/unique.ptr.observers/op_subscript.runtime.pass.cpp ---- a/libcxx/test/std/utilities/smartptr/unique.ptr/unique.ptr.class/unique.ptr.observers/op_subscript.runtime.pass.cpp -+++ b/libcxx/test/std/utilities/smartptr/unique.ptr/unique.ptr.class/unique.ptr.observers/op_subscript.runtime.pass.cpp -@@ -46,6 +46,11 @@ - template - struct CustomDeleter : std::default_delete {}; +-void elf::setSystemZTargetInfo(Ctx &ctx) { ctx.target.reset(new SystemZ(ctx)); } ++TargetInfo *elf::getSystemZTargetInfo(Ctx &ctx) { ++ static SystemZ t(ctx); ++ return &t; ++} +diff -ruN --strip-trailing-cr a/lld/ELF/Arch/X86_64.cpp b/lld/ELF/Arch/X86_64.cpp +--- a/lld/ELF/Arch/X86_64.cpp ++++ b/lld/ELF/Arch/X86_64.cpp +@@ -1237,17 +1237,21 @@ + write32le(buf + 8, ctx.in.plt->getVA() - pltEntryAddr - 12); + } -+struct NoopDeleter { -+ template -+ TEST_CONSTEXPR_CXX23 void operator()(T*) const {} -+}; -+ - TEST_CONSTEXPR_CXX23 bool test() { - // Basic test - { -@@ -112,12 +117,33 @@ - WithNonTrivialDtor<16>, - WithNonTrivialDtor<256>>; - types::for_each(TrickyCookieTypes(), [] { -- types::for_each(types::type_list, CustomDeleter>(), [] { -- std::unique_ptr p(new T[3]); -+ // Array allocated with `new T[n]`, default deleter -+ { -+ std::unique_ptr> p(new T[3]); -+ assert(p[0] == T()); -+ assert(p[1] == T()); -+ assert(p[2] == T()); -+ } -+ -+ // Array allocated with `new T[n]`, custom deleter -+ { -+ std::unique_ptr> p(new T[3]); -+ assert(p[0] == T()); -+ assert(p[1] == T()); -+ assert(p[2] == T()); -+ } -+ -+ // Array not allocated with `new T[n]`, custom deleter -+ // -+ // This test aims to ensure that the implementation doesn't try to use an array cookie -+ // when there is none. -+ { -+ T array[50] = {}; -+ std::unique_ptr p(&array[0]); - assert(p[0] == T()); - assert(p[1] == T()); - assert(p[2] == T()); -- }); -+ } - }); +-void elf::setX86_64TargetInfo(Ctx &ctx) { ++TargetInfo *elf::getX86_64TargetInfo(Ctx &ctx) { + if (ctx.arg.zRetpolineplt) { +- if (ctx.arg.zNow) +- ctx.target.reset(new RetpolineZNow(ctx)); +- else +- ctx.target.reset(new Retpoline(ctx)); +- return; ++ if (ctx.arg.zNow) { ++ static RetpolineZNow t(ctx); ++ return &t; ++ } ++ static Retpoline t(ctx); ++ return &t; } - #endif // C++20 -diff -ruN --strip-trailing-cr a/lldb/test/API/functionalities/data-formatter/data-formatter-stl/libcxx-simulators/string/main.cpp b/lldb/test/API/functionalities/data-formatter/data-formatter-stl/libcxx-simulators/string/main.cpp ---- a/lldb/test/API/functionalities/data-formatter/data-formatter-stl/libcxx-simulators/string/main.cpp -+++ b/lldb/test/API/functionalities/data-formatter/data-formatter-stl/libcxx-simulators/string/main.cpp -@@ -20,11 +20,7 @@ - // Pre-D128285 layout. - #define PACKED_ANON_STRUCT - #endif --#if REVISION <= 4 --// Pre-2a1ef74 layout. --#define NON_STANDARD_PADDING --#endif --// REVISION == 5: current layout -+// REVISION == 4: current layout - #ifdef PACKED_ANON_STRUCT - #define BEGIN_PACKED_ANON_STRUCT struct __attribute__((packed)) { -@@ -38,7 +34,6 @@ - namespace std { - namespace __lldb { +- if (ctx.arg.andFeatures & GNU_PROPERTY_X86_FEATURE_1_IBT) +- ctx.target.reset(new IntelIBT(ctx)); +- else +- ctx.target.reset(new X86_64(ctx)); ++ if (ctx.arg.andFeatures & GNU_PROPERTY_X86_FEATURE_1_IBT) { ++ static IntelIBT t(ctx); ++ return &t; ++ } ++ ++ static X86_64 t(ctx); ++ return &t; + } +diff -ruN --strip-trailing-cr a/lld/ELF/Arch/X86.cpp b/lld/ELF/Arch/X86.cpp +--- a/lld/ELF/Arch/X86.cpp ++++ b/lld/ELF/Arch/X86.cpp +@@ -706,17 +706,21 @@ + write32le(buf + 22, -off - 26); + } --#ifdef NON_STANDARD_PADDING - #if defined(ALTERNATE_LAYOUT) && defined(SUBCLASS_PADDING) - template struct __padding { - unsigned char __xx[sizeof(_CharT) - 1]; -@@ -46,13 +41,6 @@ +-void elf::setX86TargetInfo(Ctx &ctx) { ++TargetInfo *elf::getX86TargetInfo(Ctx &ctx) { + if (ctx.arg.zRetpolineplt) { +- if (ctx.arg.isPic) +- ctx.target.reset(new RetpolinePic(ctx)); +- else +- ctx.target.reset(new RetpolineNoPic(ctx)); +- return; ++ if (ctx.arg.isPic) { ++ static RetpolinePic t(ctx); ++ return &t; ++ } ++ static RetpolineNoPic t(ctx); ++ return &t; + } - template struct __padding<_CharT, 1> {}; - #endif --#else // !NON_STANDARD_PADDING --template struct __padding { -- char __padding_[_PaddingSize]; --}; -- --template <> struct __padding<0> {}; --#endif +- if (ctx.arg.andFeatures & GNU_PROPERTY_X86_FEATURE_1_IBT) +- ctx.target.reset(new IntelIBT(ctx)); +- else +- ctx.target.reset(new X86(ctx)); ++ if (ctx.arg.andFeatures & GNU_PROPERTY_X86_FEATURE_1_IBT) { ++ static IntelIBT t(ctx); ++ return &t; ++ } ++ ++ static X86 t(ctx); ++ return &t; + } +diff -ruN --strip-trailing-cr a/lld/ELF/Config.h b/lld/ELF/Config.h +--- a/lld/ELF/Config.h ++++ b/lld/ELF/Config.h +@@ -545,7 +545,7 @@ + Config arg; + LinkerDriver driver; + LinkerScript *script; +- std::unique_ptr target; ++ TargetInfo *target; + + // These variables are initialized by Writer and should not be used before + // Writer is initialized. +diff -ruN --strip-trailing-cr a/lld/ELF/Driver.cpp b/lld/ELF/Driver.cpp +--- a/lld/ELF/Driver.cpp ++++ b/lld/ELF/Driver.cpp +@@ -99,7 +99,7 @@ + driver.~LinkerDriver(); + new (&driver) LinkerDriver(*this); + script = nullptr; +- target.reset(); ++ target = nullptr; - template class basic_string { - public: -@@ -89,12 +77,7 @@ - }; - #else // !SUBCLASS_PADDING + bufferStart = nullptr; + mainPart = nullptr; +@@ -3126,7 +3126,7 @@ + // The Target instance handles target-specific stuff, such as applying + // relocations or writing a PLT section. It also contains target-dependent + // values such as a default image base address. +- setTarget(ctx); ++ ctx.target = getTarget(ctx); --#ifdef NON_STANDARD_PADDING - unsigned char __padding[sizeof(value_type) - 1]; --#else -- [[no_unique_address]] __padding __padding_; --#endif -- - #ifdef BITMASKS - unsigned char __size_; - #else // !BITMASKS -@@ -146,26 +129,21 @@ - union { - #ifdef BITMASKS - unsigned char __size_; --#else // !BITMASKS -+#else - struct { - unsigned char __is_long_ : 1; - unsigned char __size_ : 7; - }; --#endif // BITMASKS -+#endif - value_type __lx; - }; --#else // !SHORT_UNION -+#else - BEGIN_PACKED_ANON_STRUCT - unsigned char __is_long_ : 1; - unsigned char __size_ : 7; - END_PACKED_ANON_STRUCT --#ifdef NON_STANDARD_PADDING -- unsigned char __padding[sizeof(value_type) - 1]; --#else // !NON_STANDARD_PADDING -- [[no_unique_address]] __padding __padding_; --#endif // NON_STANDARD_PADDING -- --#endif // SHORT_UNION -+ char __padding_[sizeof(value_type) - 1]; -+#endif - value_type __data_[__min_cap]; - }; + ctx.arg.eflags = ctx.target->calcEFlags(); + // maxPageSize (sometimes called abi page size) is the maximum page size that +diff -ruN --strip-trailing-cr a/lld/ELF/Target.cpp b/lld/ELF/Target.cpp +--- a/lld/ELF/Target.cpp ++++ b/lld/ELF/Target.cpp +@@ -45,39 +45,39 @@ + return std::string(s); + } -diff -ruN --strip-trailing-cr a/lldb/test/API/functionalities/data-formatter/data-formatter-stl/libcxx-simulators/string/TestDataFormatterLibcxxStringSimulator.py b/lldb/test/API/functionalities/data-formatter/data-formatter-stl/libcxx-simulators/string/TestDataFormatterLibcxxStringSimulator.py ---- a/lldb/test/API/functionalities/data-formatter/data-formatter-stl/libcxx-simulators/string/TestDataFormatterLibcxxStringSimulator.py -+++ b/lldb/test/API/functionalities/data-formatter/data-formatter-stl/libcxx-simulators/string/TestDataFormatterLibcxxStringSimulator.py -@@ -27,7 +27,7 @@ +-void elf::setTarget(Ctx &ctx) { ++TargetInfo *elf::getTarget(Ctx &ctx) { + switch (ctx.arg.emachine) { + case EM_386: + case EM_IAMCU: +- return setX86TargetInfo(ctx); ++ return getX86TargetInfo(ctx); + case EM_AARCH64: +- return setAArch64TargetInfo(ctx); ++ return getAArch64TargetInfo(ctx); + case EM_AMDGPU: +- return setAMDGPUTargetInfo(ctx); ++ return getAMDGPUTargetInfo(ctx); + case EM_ARM: +- return setARMTargetInfo(ctx); ++ return getARMTargetInfo(ctx); + case EM_AVR: +- return setAVRTargetInfo(ctx); ++ return getAVRTargetInfo(ctx); + case EM_HEXAGON: +- return setHexagonTargetInfo(ctx); ++ return getHexagonTargetInfo(ctx); + case EM_LOONGARCH: +- return setLoongArchTargetInfo(ctx); ++ return getLoongArchTargetInfo(ctx); + case EM_MIPS: +- return setMipsTargetInfo(ctx); ++ return getMipsTargetInfo(ctx); + case EM_MSP430: +- return setMSP430TargetInfo(ctx); ++ return getMSP430TargetInfo(ctx); + case EM_PPC: +- return setPPCTargetInfo(ctx); ++ return getPPCTargetInfo(ctx); + case EM_PPC64: +- return setPPC64TargetInfo(ctx); ++ return getPPC64TargetInfo(ctx); + case EM_RISCV: +- return setRISCVTargetInfo(ctx); ++ return getRISCVTargetInfo(ctx); + case EM_SPARCV9: +- return setSPARCV9TargetInfo(ctx); ++ return getSPARCV9TargetInfo(ctx); + case EM_S390: +- return setSystemZTargetInfo(ctx); ++ return getSystemZTargetInfo(ctx); + case EM_X86_64: +- return setX86_64TargetInfo(ctx); ++ return getX86_64TargetInfo(ctx); + default: + fatal("unsupported e_machine value: " + Twine(ctx.arg.emachine)); + } +diff -ruN --strip-trailing-cr a/lld/ELF/Target.h b/lld/ELF/Target.h +--- a/lld/ELF/Target.h ++++ b/lld/ELF/Target.h +@@ -179,21 +179,21 @@ + uint64_t defaultImageBase = 0x10000; + }; +-void setAArch64TargetInfo(Ctx &); +-void setAMDGPUTargetInfo(Ctx &); +-void setARMTargetInfo(Ctx &); +-void setAVRTargetInfo(Ctx &); +-void setHexagonTargetInfo(Ctx &); +-void setLoongArchTargetInfo(Ctx &); +-void setMSP430TargetInfo(Ctx &); +-void setMipsTargetInfo(Ctx &); +-void setPPC64TargetInfo(Ctx &); +-void setPPCTargetInfo(Ctx &); +-void setRISCVTargetInfo(Ctx &); +-void setSPARCV9TargetInfo(Ctx &); +-void setSystemZTargetInfo(Ctx &); +-void setX86TargetInfo(Ctx &); +-void setX86_64TargetInfo(Ctx &); ++TargetInfo *getAArch64TargetInfo(Ctx &); ++TargetInfo *getAMDGPUTargetInfo(Ctx &); ++TargetInfo *getARMTargetInfo(Ctx &); ++TargetInfo *getAVRTargetInfo(Ctx &); ++TargetInfo *getHexagonTargetInfo(Ctx &); ++TargetInfo *getLoongArchTargetInfo(Ctx &); ++TargetInfo *getMSP430TargetInfo(Ctx &); ++TargetInfo *getMipsTargetInfo(Ctx &); ++TargetInfo *getPPC64TargetInfo(Ctx &); ++TargetInfo *getPPCTargetInfo(Ctx &); ++TargetInfo *getRISCVTargetInfo(Ctx &); ++TargetInfo *getSPARCV9TargetInfo(Ctx &); ++TargetInfo *getSystemZTargetInfo(Ctx &); ++TargetInfo *getX86TargetInfo(Ctx &); ++TargetInfo *getX86_64TargetInfo(Ctx &); - for v in [None, "ALTERNATE_LAYOUT"]: -- for r in range(6): -+ for r in range(5): - for c in range(3): - name = "test_r%d_c%d" % (r, c) - defines = ["REVISION=%d" % r, "COMPRESSED_PAIR_REV=%d" % c] -diff -ruN --strip-trailing-cr a/utils/bazel/llvm-project-overlay/libc/BUILD.bazel b/utils/bazel/llvm-project-overlay/libc/BUILD.bazel ---- a/utils/bazel/llvm-project-overlay/libc/BUILD.bazel -+++ b/utils/bazel/llvm-project-overlay/libc/BUILD.bazel -@@ -148,6 +148,9 @@ - libc_support_library( - name = "hdr_stdio_macros", - hdrs = ["hdr/stdio_macros.h"], -+ deps = [ -+ ":hdr_stdio_overlay", -+ ], - ) + struct ErrorPlace { + InputSectionBase *isec; +@@ -251,7 +251,7 @@ + void createTaggedSymbols(Ctx &); + void initSymbolAnchors(Ctx &); - libc_support_library( -@@ -215,11 +218,17 @@ - libc_support_library( - name = "types_off_t", - hdrs = ["hdr/types/off_t.h"], -+ deps = [ -+ ":hdr_stdio_overlay", -+ ], - ) +-void setTarget(Ctx &); ++TargetInfo *getTarget(Ctx &); - libc_support_library( - name = "types_FILE", - hdrs = ["hdr/types/FILE.h"], -+ deps = [ -+ ":hdr_stdio_overlay", -+ ], - ) + template bool isMipsPIC(const Defined *sym); - libc_support_library( diff --git a/third_party/tsl/third_party/llvm/workspace.bzl b/third_party/tsl/third_party/llvm/workspace.bzl index 50e3160cadca3..e939a2ea30102 100644 --- a/third_party/tsl/third_party/llvm/workspace.bzl +++ b/third_party/tsl/third_party/llvm/workspace.bzl @@ -4,8 +4,8 @@ load("//third_party:repo.bzl", "tf_http_archive") def repo(name): """Imports LLVM.""" - LLVM_COMMIT = "82f5acfbec65e1a645d902f746253eeaf0bd2d70" - LLVM_SHA256 = "c621acf6f202cfcbd0c73f381272b70c51c95f0ecb62c17d2d5fe0156a80082a" + LLVM_COMMIT = "2918e779a9545a66c0031b03b3af5bf4d8517cec" + LLVM_SHA256 = "efd6ac33aba00a6eac0047533a99d6d766dc9dae9ec5ba8e74ff297bfd4c9f0e" tf_http_archive( name = name,