diff options
| author | Alex Richardson <alexrichardson@google.com> | 2024-02-27 13:25:49 -0800 |
|---|---|---|
| committer | Alex Richardson <alexrichardson@google.com> | 2024-02-27 13:25:49 -0800 |
| commit | 413077b8f719e969aba07bc24b5e631f3483842b (patch) | |
| tree | 6bdd90085053dfd72ca5a480ee46e6d358337ac0 | |
| parent | eba1687159247f1dcc185a68d7ebd67e49796b09 (diff) | |
| parent | d82e93e7f129d9e8b72570efdf4a15d6ec3d4336 (diff) | |
[𝘀𝗽𝗿] changes introduced through rebaseupstream/users/arichardson/spr/main.compiler-rt-allow-running-tests-without-installing-first
Created using spr 1.3.4
[skip ci]
76 files changed, 1229 insertions, 523 deletions
diff --git a/clang/docs/ReleaseNotes.rst b/clang/docs/ReleaseNotes.rst index c60c5682dbd8..7e16b9f0c67d 100644 --- a/clang/docs/ReleaseNotes.rst +++ b/clang/docs/ReleaseNotes.rst @@ -288,6 +288,8 @@ Bug Fixes to C++ Support templates when determining the primary template of an explicit specialization. - Fixed a crash in Microsoft compatibility mode where unqualified dependent base class lookup searches the bases of an incomplete class. +- Fix a crash when an unresolved overload set is encountered on the RHS of a ``.*`` operator. + (`#53815 <https://github.com/llvm/llvm-project/issues/53815>`_) Bug Fixes to AST Handling ^^^^^^^^^^^^^^^^^^^^^^^^^ diff --git a/clang/include/clang/Basic/TargetOSMacros.def b/clang/include/clang/Basic/TargetOSMacros.def index dfc2e033f6fd..58dce330f9c8 100644 --- a/clang/include/clang/Basic/TargetOSMacros.def +++ b/clang/include/clang/Basic/TargetOSMacros.def @@ -34,18 +34,19 @@ TARGET_OS(TARGET_OS_UNIX, Triple.isOSNetBSD() || TARGET_OS(TARGET_OS_MAC, Triple.isOSDarwin()) TARGET_OS(TARGET_OS_OSX, Triple.isMacOSX()) TARGET_OS(TARGET_OS_IPHONE, Triple.isiOS() || Triple.isTvOS() || - Triple.isWatchOS()) + Triple.isWatchOS() || Triple.isXROS()) // Triple::isiOS() also includes tvOS TARGET_OS(TARGET_OS_IOS, Triple.getOS() == llvm::Triple::IOS) TARGET_OS(TARGET_OS_TV, Triple.isTvOS()) TARGET_OS(TARGET_OS_WATCH, Triple.isWatchOS()) +TARGET_OS(TARGET_OS_VISION, Triple.isXROS()) TARGET_OS(TARGET_OS_DRIVERKIT, Triple.isDriverKit()) TARGET_OS(TARGET_OS_MACCATALYST, Triple.isMacCatalystEnvironment()) TARGET_OS(TARGET_OS_SIMULATOR, Triple.isSimulatorEnvironment()) // Deprecated Apple target conditionals. TARGET_OS(TARGET_OS_EMBEDDED, (Triple.isiOS() || Triple.isTvOS() \ - || Triple.isWatchOS()) \ + || Triple.isWatchOS() || Triple.isXROS()) \ && !Triple.isMacCatalystEnvironment() \ && !Triple.isSimulatorEnvironment()) TARGET_OS(TARGET_OS_NANO, Triple.isWatchOS()) diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 1550b000a89a..d16d12fac8b0 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -397,8 +397,8 @@ CodeGenModule::CodeGenModule(ASTContext &C, // Enable TBAA unless it's suppressed. ThreadSanitizer needs TBAA even at O0. if (LangOpts.Sanitize.has(SanitizerKind::Thread) || (!CodeGenOpts.RelaxedAliasing && CodeGenOpts.OptimizationLevel > 0)) - TBAA.reset(new CodeGenTBAA(Context, TheModule, CodeGenOpts, getLangOpts(), - getCXXABI().getMangleContext())); + TBAA.reset(new CodeGenTBAA(Context, getTypes(), TheModule, CodeGenOpts, + getLangOpts(), getCXXABI().getMangleContext())); // If debug info or coverage generation is enabled, create the CGDebugInfo // object. diff --git a/clang/lib/CodeGen/CodeGenTBAA.cpp b/clang/lib/CodeGen/CodeGenTBAA.cpp index dc288bc3f615..8a0816121939 100644 --- a/clang/lib/CodeGen/CodeGenTBAA.cpp +++ b/clang/lib/CodeGen/CodeGenTBAA.cpp @@ -15,6 +15,8 @@ //===----------------------------------------------------------------------===// #include "CodeGenTBAA.h" +#include "CGRecordLayout.h" +#include "CodeGenTypes.h" #include "clang/AST/ASTContext.h" #include "clang/AST/Attr.h" #include "clang/AST/Mangle.h" @@ -26,16 +28,16 @@ #include "llvm/IR/Metadata.h" #include "llvm/IR/Module.h" #include "llvm/IR/Type.h" +#include "llvm/Support/Debug.h" using namespace clang; using namespace CodeGen; -CodeGenTBAA::CodeGenTBAA(ASTContext &Ctx, llvm::Module &M, - const CodeGenOptions &CGO, +CodeGenTBAA::CodeGenTBAA(ASTContext &Ctx, CodeGenTypes &CGTypes, + llvm::Module &M, const CodeGenOptions &CGO, const LangOptions &Features, MangleContext &MContext) - : Context(Ctx), Module(M), CodeGenOpts(CGO), - Features(Features), MContext(MContext), MDHelper(M.getContext()), - Root(nullptr), Char(nullptr) -{} + : Context(Ctx), CGTypes(CGTypes), Module(M), CodeGenOpts(CGO), + Features(Features), MContext(MContext), MDHelper(M.getContext()), + Root(nullptr), Char(nullptr) {} CodeGenTBAA::~CodeGenTBAA() { } @@ -294,14 +296,34 @@ CodeGenTBAA::CollectFields(uint64_t BaseOffset, return false; const ASTRecordLayout &Layout = Context.getASTRecordLayout(RD); + const CGRecordLayout &CGRL = CGTypes.getCGRecordLayout(RD); unsigned idx = 0; - for (RecordDecl::field_iterator i = RD->field_begin(), - e = RD->field_end(); i != e; ++i, ++idx) { - if ((*i)->isZeroSize(Context) || (*i)->isUnnamedBitfield()) + for (RecordDecl::field_iterator i = RD->field_begin(), e = RD->field_end(); + i != e; ++i, ++idx) { + if ((*i)->isZeroSize(Context)) continue; - uint64_t Offset = BaseOffset + - Layout.getFieldOffset(idx) / Context.getCharWidth(); + + uint64_t Offset = + BaseOffset + Layout.getFieldOffset(idx) / Context.getCharWidth(); + + // Create a single field for consecutive named bitfields using char as + // base type. + if ((*i)->isBitField()) { + const CGBitFieldInfo &Info = CGRL.getBitFieldInfo(*i); + if (Info.Offset != 0) + continue; + unsigned CurrentBitFieldSize = Info.StorageSize; + uint64_t Size = + llvm::divideCeil(CurrentBitFieldSize, Context.getCharWidth()); + llvm::MDNode *TBAAType = getChar(); + llvm::MDNode *TBAATag = + getAccessTagInfo(TBAAAccessInfo(TBAAType, Size)); + Fields.push_back( + llvm::MDBuilder::TBAAStructField(Offset, Size, TBAATag)); + continue; + } + QualType FieldQTy = i->getType(); if (!CollectFields(Offset, FieldQTy, Fields, MayAlias || TypeHasMayAlias(FieldQTy))) diff --git a/clang/lib/CodeGen/CodeGenTBAA.h b/clang/lib/CodeGen/CodeGenTBAA.h index a65963596fe9..aa6da2731a41 100644 --- a/clang/lib/CodeGen/CodeGenTBAA.h +++ b/clang/lib/CodeGen/CodeGenTBAA.h @@ -29,6 +29,7 @@ namespace clang { class Type; namespace CodeGen { +class CodeGenTypes; // TBAAAccessKind - A kind of TBAA memory access descriptor. enum class TBAAAccessKind : unsigned { @@ -115,6 +116,7 @@ struct TBAAAccessInfo { /// while lowering AST types to LLVM types. class CodeGenTBAA { ASTContext &Context; + CodeGenTypes &CGTypes; llvm::Module &Module; const CodeGenOptions &CodeGenOpts; const LangOptions &Features; @@ -167,8 +169,9 @@ class CodeGenTBAA { llvm::MDNode *getBaseTypeInfoHelper(const Type *Ty); public: - CodeGenTBAA(ASTContext &Ctx, llvm::Module &M, const CodeGenOptions &CGO, - const LangOptions &Features, MangleContext &MContext); + CodeGenTBAA(ASTContext &Ctx, CodeGenTypes &CGTypes, llvm::Module &M, + const CodeGenOptions &CGO, const LangOptions &Features, + MangleContext &MContext); ~CodeGenTBAA(); /// getTypeInfo - Get metadata used to describe accesses to objects of the diff --git a/clang/lib/Sema/SemaOverload.cpp b/clang/lib/Sema/SemaOverload.cpp index ecad2b968165..7d38043890ca 100644 --- a/clang/lib/Sema/SemaOverload.cpp +++ b/clang/lib/Sema/SemaOverload.cpp @@ -14571,6 +14571,23 @@ ExprResult Sema::CreateOverloadedBinOp(SourceLocation OpLoc, CurFPFeatureOverrides()); } + // If this is the .* operator, which is not overloadable, just + // create a built-in binary operator. + if (Opc == BO_PtrMemD) { + auto CheckPlaceholder = [&](Expr *&Arg) { + ExprResult Res = CheckPlaceholderExpr(Arg); + if (Res.isUsable()) + Arg = Res.get(); + return !Res.isUsable(); + }; + + // CreateBuiltinBinOp() doesn't like it if we tell it to create a '.*' + // expression that contains placeholders (in either the LHS or RHS). + if (CheckPlaceholder(Args[0]) || CheckPlaceholder(Args[1])) + return ExprError(); + return CreateBuiltinBinOp(OpLoc, Opc, Args[0], Args[1]); + } + // Always do placeholder-like conversions on the RHS. if (checkPlaceholderForOverload(*this, Args[1])) return ExprError(); @@ -14590,11 +14607,6 @@ ExprResult Sema::CreateOverloadedBinOp(SourceLocation OpLoc, if (Opc == BO_Assign && !Args[0]->getType()->isOverloadableType()) return CreateBuiltinBinOp(OpLoc, Opc, Args[0], Args[1]); - // If this is the .* operator, which is not overloadable, just - // create a built-in binary operator. - if (Opc == BO_PtrMemD) - return CreateBuiltinBinOp(OpLoc, Opc, Args[0], Args[1]); - // Build the overload set. OverloadCandidateSet CandidateSet(OpLoc, OverloadCandidateSet::CSK_Operator, OverloadCandidateSet::OperatorRewriteInfo( diff --git a/clang/test/CodeGen/tbaa-struct.cpp b/clang/test/CodeGen/tbaa-struct.cpp index 28c7d396121a..883c982be26c 100644 --- a/clang/test/CodeGen/tbaa-struct.cpp +++ b/clang/test/CodeGen/tbaa-struct.cpp @@ -162,11 +162,11 @@ void copy10(NamedBitfields3 *a1, NamedBitfields3 *a2) { // CHECK-OLD: [[TS3]] = !{i64 0, i64 8, !{{.*}}, i64 0, i64 2, !{{.*}}, i64 4, i64 8, !{{.*}}} // CHECK-OLD: [[TS4]] = !{i64 0, i64 1, [[TAG_CHAR]], i64 1, i64 1, [[TAG_CHAR]], i64 2, i64 1, [[TAG_CHAR]]} // CHECK-OLD: [[TS5]] = !{i64 0, i64 1, [[TAG_CHAR]], i64 4, i64 1, [[TAG_CHAR]], i64 5, i64 1, [[TAG_CHAR]]} -// CHECK-OLD: [[TS6]] = !{i64 0, i64 4, [[TAG_INT]], i64 1, i64 4, [[TAG_INT]], i64 2, i64 1, [[TAG_CHAR]], i64 8, i64 8, [[TAG_DOUBLE:!.+]]} +// CHECK-OLD: [[TS6]] = !{i64 0, i64 2, [[TAG_CHAR]], i64 2, i64 1, [[TAG_CHAR]], i64 8, i64 8, [[TAG_DOUBLE:!.+]]} // CHECK-OLD: [[TAG_DOUBLE]] = !{[[DOUBLE:!.+]], [[DOUBLE]], i64 0} // CHECK-OLD [[DOUBLE]] = !{!"double", [[CHAR]], i64 0} -// CHECK-OLD: [[TS7]] = !{i64 0, i64 1, [[TAG_CHAR]], i64 1, i64 1, [[TAG_CHAR]], i64 2, i64 1, [[TAG_CHAR]], i64 3, i64 4, [[TAG_INT]], i64 3, i64 4, [[TAG_INT]], i64 4, i64 1, [[TAG_CHAR]], i64 8, i64 8, [[TAG_DOUBLE]], i64 16, i64 4, [[TAG_INT]]} -// CHECK-OLD: [[TS8]] = !{i64 1, i64 4, [[TAG_INT]], i64 2, i64 4, [[TAG_INT]], i64 8, i64 8, [[TAG_DOUBLE]]} +// CHECK-OLD: [[TS7]] = !{i64 0, i64 1, [[TAG_CHAR]], i64 1, i64 1, [[TAG_CHAR]], i64 2, i64 1, [[TAG_CHAR]], i64 3, i64 1, [[TAG_CHAR]], i64 4, i64 1, [[TAG_CHAR]], i64 8, i64 8, [[TAG_DOUBLE]], i64 16, i64 1, [[TAG_CHAR]]} +// CHECK-OLD: [[TS8]] = !{i64 0, i64 4, [[TAG_CHAR]], i64 8, i64 8, [[TAG_DOUBLE]]} // CHECK-NEW-DAG: [[TYPE_char:!.*]] = !{{{.*}}, i64 1, !"omnipotent char"} // CHECK-NEW-DAG: [[TAG_char]] = !{[[TYPE_char]], [[TYPE_char]], i64 0, i64 0} diff --git a/clang/test/Driver/fdefine-target-os-macros.c b/clang/test/Driver/fdefine-target-os-macros.c index d7379dd3d539..a4de51e8e724 100644 --- a/clang/test/Driver/fdefine-target-os-macros.c +++ b/clang/test/Driver/fdefine-target-os-macros.c @@ -12,6 +12,7 @@ // RUN: -DIOS=0 \ // RUN: -DTV=0 \ // RUN: -DWATCH=0 \ +// RUN: -DVISION=0 \ // RUN: -DDRIVERKIT=0 \ // RUN: -DMACCATALYST=0 \ // RUN: -DEMBEDDED=0 \ @@ -27,6 +28,7 @@ // RUN: -DIOS=1 \ // RUN: -DTV=0 \ // RUN: -DWATCH=0 \ +// RUN: -DVISION=0 \ // RUN: -DDRIVERKIT=0 \ // RUN: -DMACCATALYST=0 \ // RUN: -DEMBEDDED=1 \ @@ -42,6 +44,7 @@ // RUN: -DIOS=1 \ // RUN: -DTV=0 \ // RUN: -DWATCH=0 \ +// RUN: -DVISION=0 \ // RUN: -DDRIVERKIT=0 \ // RUN: -DMACCATALYST=1 \ // RUN: -DEMBEDDED=0 \ @@ -57,6 +60,7 @@ // RUN: -DIOS=1 \ // RUN: -DTV=0 \ // RUN: -DWATCH=0 \ +// RUN: -DVISION=0 \ // RUN: -DDRIVERKIT=0 \ // RUN: -DMACCATALYST=0 \ // RUN: -DEMBEDDED=0 \ @@ -72,6 +76,7 @@ // RUN: -DIOS=0 \ // RUN: -DTV=1 \ // RUN: -DWATCH=0 \ +// RUN: -DVISION=0 \ // RUN: -DDRIVERKIT=0 \ // RUN: -DMACCATALYST=0 \ // RUN: -DEMBEDDED=1 \ @@ -87,6 +92,7 @@ // RUN: -DIOS=0 \ // RUN: -DTV=1 \ // RUN: -DWATCH=0 \ +// RUN: -DVISION=0 \ // RUN: -DDRIVERKIT=0 \ // RUN: -DMACCATALYST=0 \ // RUN: -DEMBEDDED=0 \ @@ -102,6 +108,7 @@ // RUN: -DIOS=0 \ // RUN: -DTV=0 \ // RUN: -DWATCH=1 \ +// RUN: -DVISION=0 \ // RUN: -DDRIVERKIT=0 \ // RUN: -DMACCATALYST=0 \ // RUN: -DEMBEDDED=1 \ @@ -117,6 +124,39 @@ // RUN: -DIOS=0 \ // RUN: -DTV=0 \ // RUN: -DWATCH=1 \ +// RUN: -DVISION=0 \ +// RUN: -DDRIVERKIT=0 \ +// RUN: -DMACCATALYST=0 \ +// RUN: -DEMBEDDED=0 \ +// RUN: -DSIMULATOR=1 \ +// RUN: -DWINDOWS=0 \ +// RUN: -DLINUX=0 \ +// RUN: -DUNIX=0 + +// RUN: %clang -dM -E --target=arm64-apple-xros %s 2>&1 \ +// RUN: | FileCheck %s -DMAC=1 \ +// RUN: -DOSX=0 \ +// RUN: -DIPHONE=1 \ +// RUN: -DIOS=0 \ +// RUN: -DTV=0 \ +// RUN: -DWATCH=0 \ +// RUN: -DVISION=1 \ +// RUN: -DDRIVERKIT=0 \ +// RUN: -DMACCATALYST=0 \ +// RUN: -DEMBEDDED=1 \ +// RUN: -DSIMULATOR=0 \ +// RUN: -DWINDOWS=0 \ +// RUN: -DLINUX=0 \ +// RUN: -DUNIX=0 + +// RUN: %clang -dM -E --target=arm64-apple-xros-simulator %s 2>&1 \ +// RUN: | FileCheck %s -DMAC=1 \ +// RUN: -DOSX=0 \ +// RUN: -DIPHONE=1 \ +// RUN: -DIOS=0 \ +// RUN: -DTV=0 \ +// RUN: -DWATCH=0 \ +// RUN: -DVISION=1 \ // RUN: -DDRIVERKIT=0 \ // RUN: -DMACCATALYST=0 \ // RUN: -DEMBEDDED=0 \ @@ -132,6 +172,7 @@ // RUN: -DIOS=0 \ // RUN: -DTV=0 \ // RUN: -DWATCH=0 \ +// RUN: -DVISION=0 \ // RUN: -DDRIVERKIT=1 \ // RUN: -DMACCATALYST=0 \ // RUN: -DEMBEDDED=0 \ @@ -148,6 +189,7 @@ // RUN: -DIOS=0 \ // RUN: -DTV=0 \ // RUN: -DWATCH=0 \ +// RUN: -DVISION=0 \ // RUN: -DDRIVERKIT=0 \ // RUN: -DMACCATALYST=0 \ // RUN: -DEMBEDDED=0 \ @@ -164,6 +206,7 @@ // RUN: -DIOS=0 \ // RUN: -DTV=0 \ // RUN: -DWATCH=0 \ +// RUN: -DVISION=0 \ // RUN: -DDRIVERKIT=0 \ // RUN: -DMACCATALYST=0 \ // RUN: -DEMBEDDED=0 \ @@ -180,6 +223,7 @@ // RUN: -DIOS=0 \ // RUN: -DTV=0 \ // RUN: -DWATCH=0 \ +// RUN: -DVISION=0 \ // RUN: -DDRIVERKIT=0 \ // RUN: -DMACCATALYST=0 \ // RUN: -DEMBEDDED=0 \ @@ -196,6 +240,7 @@ // RUN: -DIOS=0 \ // RUN: -DTV=0 \ // RUN: -DWATCH=0 \ +// RUN: -DVISION=0 \ // RUN: -DDRIVERKIT=0 \ // RUN: -DMACCATALYST=0 \ // RUN: -DEMBEDDED=0 \ @@ -226,6 +271,7 @@ // CHECK-DAG: #define TARGET_OS_IOS [[IOS]] // CHECK-DAG: #define TARGET_OS_TV [[TV]] // CHECK-DAG: #define TARGET_OS_WATCH [[WATCH]] +// CHECK-DAG: #define TARGET_OS_VISION [[VISION]] // CHECK-DAG: #define TARGET_OS_DRIVERKIT [[DRIVERKIT]] // CHECK-DAG: #define TARGET_OS_MACCATALYST [[MACCATALYST]] // CHECK-DAG: #define TARGET_OS_SIMULATOR [[SIMULATOR]] diff --git a/clang/test/SemaCXX/gh53815.cpp b/clang/test/SemaCXX/gh53815.cpp new file mode 100644 index 000000000000..326c911c7bfa --- /dev/null +++ b/clang/test/SemaCXX/gh53815.cpp @@ -0,0 +1,21 @@ +// RUN: %clang_cc1 -fsyntax-only -verify -std=c++20 %s +// expected-no-diagnostics + +// Check that we don't crash due to forgetting to check for placeholders +// in the RHS of '.*'. + +template <typename Fn> +static bool has_explicitly_named_overload() { + return requires { Fn().*&Fn::operator(); }; +} + +int main() { + has_explicitly_named_overload<decltype([](auto){})>(); +} + +template <typename Fn> +constexpr bool has_explicitly_named_overload_2() { + return requires { Fn().*&Fn::operator(); }; +} + +static_assert(!has_explicitly_named_overload_2<decltype([](auto){})>()); diff --git a/compiler-rt/CMakeLists.txt b/compiler-rt/CMakeLists.txt index bbb4e8d7c333..8a2b138d8d70 100644 --- a/compiler-rt/CMakeLists.txt +++ b/compiler-rt/CMakeLists.txt @@ -771,8 +771,6 @@ mark_as_advanced(COMPILER_RT_ENABLE_INTERNAL_SYMBOLIZER) add_subdirectory(lib) if(COMPILER_RT_INCLUDE_TESTS) - add_subdirectory(unittests) - add_subdirectory(test) # Don't build llvm-lit for runtimes-build, it will clean up map_config. if (COMPILER_RT_STANDALONE_BUILD AND NOT LLVM_RUNTIMES_BUILD) # If we have a valid source tree, generate llvm-lit into the bin directory. @@ -782,11 +780,17 @@ if(COMPILER_RT_INCLUDE_TESTS) # Needed for lit support in standalone builds. include(AddLLVM) add_subdirectory(${LLVM_MAIN_SRC_DIR}/utils/llvm-lit ${CMAKE_CURRENT_BINARY_DIR}/llvm-lit) + # Ensure that the testsuite uses the local lit rather than + # LLVM_INSTALL_DIR/bin/llvm-lit (which probably does not exist). + get_llvm_lit_path(_base_dir _file_name) + set(LLVM_EXTERNAL_LIT "${_base_dir}/${_file_name}" CACHE STRING "Command used to spawn lit" FORCE) elseif(NOT EXISTS ${LLVM_EXTERNAL_LIT}) message(WARNING "Could not find LLVM source directory and LLVM_EXTERNAL_LIT does not" "point to a valid file. You will not be able to run tests.") endif() endif() + add_subdirectory(unittests) + add_subdirectory(test) endif() add_subdirectory(tools) diff --git a/compiler-rt/cmake/Modules/CompilerRTCompile.cmake b/compiler-rt/cmake/Modules/CompilerRTCompile.cmake index 2bf115973a49..64e7acb9afd8 100644 --- a/compiler-rt/cmake/Modules/CompilerRTCompile.cmake +++ b/compiler-rt/cmake/Modules/CompilerRTCompile.cmake @@ -70,14 +70,9 @@ function(clang_compile object_file source) if (TARGET CompilerRTUnitTestCheckCxx) list(APPEND SOURCE_DEPS CompilerRTUnitTestCheckCxx) endif() - string(REGEX MATCH "[.](cc|cpp)$" is_cxx ${source_rpath}) - if (is_cxx) - set(compiler ${COMPILER_RT_TEST_COMPILER}) - else() - set(compiler ${COMPILER_RT_TEST_CXX_COMPILER}) - endif() if(COMPILER_RT_STANDALONE_BUILD) # Only add global flags in standalone build. + string(REGEX MATCH "[.](cc|cpp)$" is_cxx ${source_rpath}) if(is_cxx) string(REPLACE " " ";" global_flags "${CMAKE_CXX_FLAGS}") else() @@ -107,7 +102,7 @@ function(clang_compile object_file source) add_custom_command( OUTPUT ${object_file} - COMMAND ${compiler} ${compile_flags} -c + COMMAND ${COMPILER_RT_TEST_COMPILER} ${compile_flags} -c -o "${object_file}" ${source_rpath} MAIN_DEPENDENCY ${source} diff --git a/compiler-rt/test/dfsan/reaches_function.c b/compiler-rt/test/dfsan/reaches_function.c index afcd08cee76a..9e2bcee935b2 100644 --- a/compiler-rt/test/dfsan/reaches_function.c +++ b/compiler-rt/test/dfsan/reaches_function.c @@ -32,7 +32,7 @@ void my_dfsan_reaches_function_callback(dfsan_label label, dfsan_origin origin, __attribute__((noinline)) uint64_t add(uint64_t *a, uint64_t *b) { return *a + *b; - // CHECK: test/dfsan/reaches_function.c:[[# @LINE - 1]] add.dfsan + // CHECK: {{.*}}compiler-rt/test/dfsan/reaches_function.c:[[# @LINE - 1]] add.dfsan // CHECK-ORIGIN-TRACKING: Origin value: 0x10000002, Taint value was stored to memory at // CHECK-ORIGIN-TRACKING: #0 {{.*}} in add.dfsan {{.*}}compiler-rt/test/dfsan/reaches_function.c:[[# @LINE - 3]]:{{.*}} // CHECK-ORIGIN-TRACKING: Origin value: 0x1, Taint value was created at @@ -54,7 +54,7 @@ int main(int argc, char *argv[]) { dfsan_set_label(8, &a, sizeof(a)); uint64_t c = add(&a, &b); - // CHECK: test/dfsan/reaches_function.c:[[# @LINE - 1]] main + // CHECK: {{.*}}compiler-rt/test/dfsan/reaches_function.c:[[# @LINE - 1]] main // CHECK-ORIGIN-TRACKING: Origin value: 0x10000002, Taint value was stored to memory at // CHECK-ORIGIN-TRACKING: #0 {{.*}} in add.dfsan {{.*}}compiler-rt/test/dfsan/reaches_function.c:{{.*}} // CHECK-ORIGIN-TRACKING: Origin value: 0x1, Taint value was created at diff --git a/flang/include/flang/Lower/PFTBuilder.h b/flang/include/flang/Lower/PFTBuilder.h index c2b0fdbf357c..9913f584133f 100644 --- a/flang/include/flang/Lower/PFTBuilder.h +++ b/flang/include/flang/Lower/PFTBuilder.h @@ -138,7 +138,8 @@ using Directives = std::tuple<parser::CompilerDirective, parser::OpenACCConstruct, parser::OpenACCRoutineConstruct, parser::OpenACCDeclarativeConstruct, parser::OpenMPConstruct, - parser::OpenMPDeclarativeConstruct, parser::OmpEndLoopDirective>; + parser::OpenMPDeclarativeConstruct, parser::OmpEndLoopDirective, + parser::CUFKernelDoConstruct>; using DeclConstructs = std::tuple<parser::OpenMPDeclarativeConstruct, parser::OpenACCDeclarativeConstruct>; @@ -178,7 +179,7 @@ static constexpr bool isNopConstructStmt{common::HasMember< template <typename A> static constexpr bool isExecutableDirective{common::HasMember< A, std::tuple<parser::CompilerDirective, parser::OpenACCConstruct, - parser::OpenMPConstruct>>}; + parser::OpenMPConstruct, parser::CUFKernelDoConstruct>>}; template <typename A> static constexpr bool isFunctionLike{common::HasMember< diff --git a/flang/include/flang/Optimizer/Dialect/FIROps.td b/flang/include/flang/Optimizer/Dialect/FIROps.td index 08239230f793..db5e5f4bc682 100644 --- a/flang/include/flang/Optimizer/Dialect/FIROps.td +++ b/flang/include/flang/Optimizer/Dialect/FIROps.td @@ -3127,4 +3127,31 @@ def fir_BoxOffsetOp : fir_Op<"box_offset", [NoMemoryEffect]> { ]; } +def fir_CUDAKernelOp : fir_Op<"cuda_kernel", [AttrSizedOperandSegments, + DeclareOpInterfaceMethods<LoopLikeOpInterface>]> { + + let arguments = (ins + Variadic<I32>:$grid, // empty means `*` + Variadic<I32>:$block, // empty means `*` + Optional<I32>:$stream, + Variadic<Index>:$lowerbound, + Variadic<Index>:$upperbound, + Variadic<Index>:$step, + OptionalAttr<I64Attr>:$n + ); + + let regions = (region AnyRegion:$region); + + let assemblyFormat = [{ + `<` `<` `<` custom<CUFKernelValues>($grid, type($grid)) `,` + custom<CUFKernelValues>($block, type($block)) + ( `,` `stream` `=` $stream^ )? `>` `>` `>` + custom<CUFKernelLoopControl>($region, $lowerbound, type($lowerbound), + $upperbound, type($upperbound), $step, type($step)) + attr-dict + }]; + + let hasVerifier = 1; +} + #endif diff --git a/flang/lib/Lower/Bridge.cpp b/flang/lib/Lower/Bridge.cpp index 83555e7cd82e..f865b53f74de 100644 --- a/flang/lib/Lower/Bridge.cpp +++ b/flang/lib/Lower/Bridge.cpp @@ -2474,6 +2474,135 @@ private: // Handled by genFIR(const Fortran::parser::OpenACCDeclarativeConstruct &) } + void genFIR(const Fortran::parser::CUFKernelDoConstruct &kernel) { + localSymbols.pushScope(); + const Fortran::parser::CUFKernelDoConstruct::Directive &dir = + std::get<Fortran::parser::CUFKernelDoConstruct::Directive>(kernel.t); + + mlir::Location loc = genLocation(dir.source); + + Fortran::lower::StatementContext stmtCtx; + + unsigned nestedLoops = 1; + + const auto &nLoops = + std::get<std::optional<Fortran::parser::ScalarIntConstantExpr>>(dir.t); + if (nLoops) + nestedLoops = *Fortran::semantics::GetIntValue(*nLoops); + + mlir::IntegerAttr n; + if (nestedLoops > 1) + n = builder->getIntegerAttr(builder->getI64Type(), nestedLoops); + + const std::list<Fortran::parser::ScalarIntExpr> &grid = std::get<1>(dir.t); + const std::list<Fortran::parser::ScalarIntExpr> &block = std::get<2>(dir.t); + const std::optional<Fortran::parser::ScalarIntExpr> &stream = + std::get<3>(dir.t); + + llvm::SmallVector<mlir::Value> gridValues; + for (const Fortran::parser::ScalarIntExpr &expr : grid) + gridValues.push_back(fir::getBase( + genExprValue(*Fortran::semantics::GetExpr(expr), stmtCtx))); + llvm::SmallVector<mlir::Value> blockValues; + for (const Fortran::parser::ScalarIntExpr &expr : block) + blockValues.push_back(fir::getBase( + genExprValue(*Fortran::semantics::GetExpr(expr), stmtCtx))); + mlir::Value streamValue; + if (stream) + streamValue = fir::getBase( + genExprValue(*Fortran::semantics::GetExpr(*stream), stmtCtx)); + + const auto &outerDoConstruct = + std::get<std::optional<Fortran::parser::DoConstruct>>(kernel.t); + + llvm::SmallVector<mlir::Location> locs; + locs.push_back(loc); + llvm::SmallVector<mlir::Value> lbs, ubs, steps; + + mlir::Type idxTy = builder->getIndexType(); + + llvm::SmallVector<mlir::Type> ivTypes; + llvm::SmallVector<mlir::Location> ivLocs; + llvm::SmallVector<mlir::Value> ivValues; + for (unsigned i = 0; i < nestedLoops; ++i) { + const Fortran::parser::LoopControl *loopControl; + Fortran::lower::pft::Evaluation *loopEval = + &getEval().getFirstNestedEvaluation(); + + mlir::Location crtLoc = loc; + if (i == 0) { + loopControl = &*outerDoConstruct->GetLoopControl(); + crtLoc = + genLocation(Fortran::parser::FindSourceLocation(outerDoConstruct)); + } else { + auto *doCons = loopEval->getIf<Fortran::parser::DoConstruct>(); + assert(doCons && "expect do construct"); + loopControl = &*doCons->GetLoopControl(); + crtLoc = genLocation(Fortran::parser::FindSourceLocation(*doCons)); + } + + locs.push_back(crtLoc); + + const Fortran::parser::LoopControl::Bounds *bounds = + std::get_if<Fortran::parser::LoopControl::Bounds>(&loopControl->u); + assert(bounds && "Expected bounds on the loop construct"); + + Fortran::semantics::Symbol &ivSym = + bounds->name.thing.symbol->GetUltimate(); + ivValues.push_back(getSymbolAddress(ivSym)); + + lbs.push_back(builder->createConvert( + crtLoc, idxTy, + fir::getBase(genExprValue(*Fortran::semantics::GetExpr(bounds->lower), + stmtCtx)))); + ubs.push_back(builder->createConvert( + crtLoc, idxTy, + fir::getBase(genExprValue(*Fortran::semantics::GetExpr(bounds->upper), + stmtCtx)))); + if (bounds->step) + steps.push_back(fir::getBase( + genExprValue(*Fortran::semantics::GetExpr(bounds->step), stmtCtx))); + else // If `step` is not present, assume it is `1`. + steps.push_back(builder->createIntegerConstant(loc, idxTy, 1)); + + ivTypes.push_back(idxTy); + ivLocs.push_back(crtLoc); + if (i < nestedLoops - 1) + loopEval = &*std::next(loopEval->getNestedEvaluations().begin()); + } + + auto op = builder->create<fir::CUDAKernelOp>( + loc, gridValues, blockValues, streamValue, lbs, ubs, steps, n); + builder->createBlock(&op.getRegion(), op.getRegion().end(), ivTypes, + ivLocs); + mlir::Block &b = op.getRegion().back(); + builder->setInsertionPointToStart(&b); + + for (auto [arg, value] : llvm::zip( + op.getLoopRegions().front()->front().getArguments(), ivValues)) { + mlir::Value convArg = + builder->createConvert(loc, fir::unwrapRefType(value.getType()), arg); + builder->create<fir::StoreOp>(loc, convArg, value); + } + + builder->create<fir::FirEndOp>(loc); + builder->setInsertionPointToStart(&b); + + Fortran::lower::pft::Evaluation *crtEval = &getEval(); + if (crtEval->lowerAsStructured()) { + crtEval = &crtEval->getFirstNestedEvaluation(); + for (int64_t i = 1; i < nestedLoops; i++) + crtEval = &*std::next(crtEval->getNestedEvaluations().begin()); + } + + // Generate loop body + for (Fortran::lower::pft::Evaluation &e : crtEval->getNestedEvaluations()) + genFIR(e); + + builder->setInsertionPointAfter(op); + localSymbols.popScope(); + } + void genFIR(const Fortran::parser::OpenMPConstruct &omp) { mlir::OpBuilder::InsertPoint insertPt = builder->saveInsertionPoint(); genOpenMPConstruct(*this, localSymbols, bridge.getSemanticsContext(), diff --git a/flang/lib/Optimizer/Dialect/FIROps.cpp b/flang/lib/Optimizer/Dialect/FIROps.cpp index 0a534cdb3c48..9bb10a42a399 100644 --- a/flang/lib/Optimizer/Dialect/FIROps.cpp +++ b/flang/lib/Optimizer/Dialect/FIROps.cpp @@ -3866,6 +3866,103 @@ mlir::LogicalResult fir::DeclareOp::verify() { return fortranVar.verifyDeclareLikeOpImpl(getMemref()); } +llvm::SmallVector<mlir::Region *> fir::CUDAKernelOp::getLoopRegions() { + return {&getRegion()}; +} + +mlir::ParseResult parseCUFKernelValues( + mlir::OpAsmParser &parser, + llvm::SmallVectorImpl<mlir::OpAsmParser::UnresolvedOperand> &values, + llvm::SmallVectorImpl<mlir::Type> &types) { + if (mlir::succeeded(parser.parseOptionalStar())) + return mlir::success(); + + if (parser.parseOptionalLParen()) { + if (mlir::failed(parser.parseCommaSeparatedList( + mlir::AsmParser::Delimiter::None, [&]() { + if (parser.parseOperand(values.emplace_back())) + return mlir::failure(); + return mlir::success(); + }))) + return mlir::failure(); + if (parser.parseRParen()) + return mlir::failure(); + } else { + if (parser.parseOperand(values.emplace_back())) + return mlir::failure(); + return mlir::success(); + } + return mlir::success(); +} + +void printCUFKernelValues(mlir::OpAsmPrinter &p, mlir::Operation *op, + mlir::ValueRange values, mlir::TypeRange types) { + if (values.empty()) + p << "*"; + + if (values.size() > 1) + p << "("; + llvm::interleaveComma(values, p, [&p](mlir::Value v) { p << v; }); + if (values.size() > 1) + p << ")"; +} + +mlir::ParseResult parseCUFKernelLoopControl( + mlir::OpAsmParser &parser, mlir::Region ®ion, + llvm::SmallVectorImpl<mlir::OpAsmParser::UnresolvedOperand> &lowerbound, + llvm::SmallVectorImpl<mlir::Type> &lowerboundType, + llvm::SmallVectorImpl<mlir::OpAsmParser::UnresolvedOperand> &upperbound, + llvm::SmallVectorImpl<mlir::Type> &upperboundType, + llvm::SmallVectorImpl<mlir::OpAsmParser::UnresolvedOperand> &step, + llvm::SmallVectorImpl<mlir::Type> &stepType) { + + llvm::SmallVector<mlir::OpAsmParser::Argument> inductionVars; + if (parser.parseLParen() || + parser.parseArgumentList(inductionVars, + mlir::OpAsmParser::Delimiter::None, + /*allowType=*/true) || + parser.parseRParen() || parser.parseEqual() || parser.parseLParen() || + parser.parseOperandList(lowerbound, inductionVars.size(), + mlir::OpAsmParser::Delimiter::None) || + parser.parseColonTypeList(lowerboundType) || parser.parseRParen() || + parser.parseKeyword("to") || parser.parseLParen() || + parser.parseOperandList(upperbound, inductionVars.size(), + mlir::OpAsmParser::Delimiter::None) || + parser.parseColonTypeList(upperboundType) || parser.parseRParen() || + parser.parseKeyword("step") || parser.parseLParen() || + parser.parseOperandList(step, inductionVars.size(), + mlir::OpAsmParser::Delimiter::None) || + parser.parseColonTypeList(stepType) || parser.parseRParen()) + return mlir::failure(); + return parser.parseRegion(region, inductionVars); +} + +void printCUFKernelLoopControl( + mlir::OpAsmPrinter &p, mlir::Operation *op, mlir::Region ®ion, + mlir::ValueRange lowerbound, mlir::TypeRange lowerboundType, + mlir::ValueRange upperbound, mlir::TypeRange upperboundType, + mlir::ValueRange steps, mlir::TypeRange stepType) { + mlir::ValueRange regionArgs = region.front().getArguments(); + if (!regionArgs.empty()) { + p << "("; + llvm::interleaveComma( + regionArgs, p, [&p](mlir::Value v) { p << v << " : " << v.getType(); }); + p << ") = (" << lowerbound << " : " << lowerboundType << ") to (" + << upperbound << " : " << upperboundType << ") " + << " step (" << steps << " : " << stepType << ") "; + } + p.printRegion(region, /*printEntryBlockArgs=*/false); +} + +mlir::LogicalResult fir::CUDAKernelOp::verify() { + if (getLowerbound().size() != getUpperbound().size() || + getLowerbound().size() != getStep().size()) + return emitOpError( + "expect same number of values in lowerbound, upperbound and step"); + + return mlir::success(); +} + //===----------------------------------------------------------------------===// // FIROpsDialect //===----------------------------------------------------------------------===// diff --git a/flang/test/Lower/CUDA/cuda-kernel-loop-directive.cuf b/flang/test/Lower/CUDA/cuda-kernel-loop-directive.cuf new file mode 100644 index 000000000000..db628fe756b9 --- /dev/null +++ b/flang/test/Lower/CUDA/cuda-kernel-loop-directive.cuf @@ -0,0 +1,51 @@ +! RUN: bbc -emit-hlfir -fcuda %s -o - | FileCheck %s + +! Test lowering of CUDA kernel loop directive. + +subroutine sub1() + integer :: i, j + integer, parameter :: n = 100 + real :: a(n), b(n) + real :: c(n,n), d(n,n) + +! CHECK-LABEL: func.func @_QPsub1() +! CHECK: %[[IV:.*]]:2 = hlfir.declare %{{.*}} {uniq_name = "_QFsub1Ei"} : (!fir.ref<i32>) -> (!fir.ref<i32>, !fir.ref<i32>) + + !$cuf kernel do <<< 1, 2 >>> + do i = 1, n + a(i) = a(i) * b(i) + end do + +! CHECK: %[[LB:.*]] = fir.convert %c1{{.*}} : (i32) -> index +! CHECK: %[[UB:.*]] = fir.convert %c100{{.*}} : (i32) -> index +! CHECK: %[[STEP:.*]] = arith.constant 1 : index +! CHECK: fir.cuda_kernel<<<%c1_i32, %c2_i32>>> (%[[ARG0:.*]] : index) = (%[[LB]] : index) to (%[[UB]] : index) step (%[[STEP]] : index) +! CHECK-NOT: fir.do_loop +! CHECK: %[[ARG0_I32:.*]] = fir.convert %[[ARG0]] : (index) -> i32 +! CHECK: fir.store %[[ARG0_I32]] to %[[IV]]#1 : !fir.ref<i32> + + + !$cuf kernel do <<< *, * >>> + do i = 1, n + a(i) = a(i) * b(i) + end do + +! CHECK: fir.cuda_kernel<<<*, *>>> (%{{.*}} : index) = (%{{.*}} : index) to (%{{.*}} : index) step (%{{.*}} : index) + + !$cuf kernel do(2) <<< 1, (256,1) >>> + do i = 1, n + do j = 1, n + c(i,j) = c(i,j) * d(i,j) + end do + end do + +! CHECK: fir.cuda_kernel<<<%c1{{.*}}, (%c256{{.*}}, %c1{{.*}})>>> (%{{.*}} : index, %{{.*}} : index) = (%{{.*}}, %{{.*}} : index, index) to (%{{.*}}, %{{.*}} : index, index) step (%{{.*}}, %{{.*}} : index, index) +! CHECK: {n = 2 : i64} + +! TODO: currently these trigger error in the parser +! !$cuf kernel do(2) <<< (1,*), (256,1) >>> +! !$cuf kernel do(2) <<< (*,*), (32,4) >>> +end + + + diff --git a/libc/cmake/modules/LLVMLibCCompileOptionRules.cmake b/libc/cmake/modules/LLVMLibCCompileOptionRules.cmake index c7ccd392354c..72b04822d8b8 100644 --- a/libc/cmake/modules/LLVMLibCCompileOptionRules.cmake +++ b/libc/cmake/modules/LLVMLibCCompileOptionRules.cmake @@ -87,6 +87,7 @@ function(_get_common_compile_options output_var flags) list(APPEND compile_options "-fvisibility=hidden") list(APPEND compile_options "-fconvergent-functions") list(APPEND compile_options "-flto") + list(APPEND compile_options "-Wno-multi-gpu") if(LIBC_TARGET_ARCHITECTURE_IS_NVPTX) list(APPEND compile_options "-Wno-unknown-cuda-version") diff --git a/libc/cmake/modules/LLVMLibCTestRules.cmake b/libc/cmake/modules/LLVMLibCTestRules.cmake index 76ce6754bd73..836e15d34741 100644 --- a/libc/cmake/modules/LLVMLibCTestRules.cmake +++ b/libc/cmake/modules/LLVMLibCTestRules.cmake @@ -463,7 +463,7 @@ function(add_integration_test test_name) if(LIBC_TARGET_ARCHITECTURE_IS_AMDGPU) target_link_options(${fq_build_target_name} PRIVATE - ${LIBC_COMPILE_OPTIONS_DEFAULT} + ${LIBC_COMPILE_OPTIONS_DEFAULT} -Wno-multi-gpu -mcpu=${LIBC_GPU_TARGET_ARCHITECTURE} -flto "-Wl,-mllvm,-amdgpu-lower-global-ctor-dtor=0" -nostdlib -static "-Wl,-mllvm,-amdhsa-code-object-version=${LIBC_GPU_CODE_OBJECT_VERSION}") @@ -471,7 +471,7 @@ function(add_integration_test test_name) # We need to use the internal object versions for NVPTX. set(internal_suffix ".__internal__") target_link_options(${fq_build_target_name} PRIVATE - ${LIBC_COMPILE_OPTIONS_DEFAULT} + ${LIBC_COMPILE_OPTIONS_DEFAULT} -Wno-multi-gpu "-Wl,--suppress-stack-size-warning" -march=${LIBC_GPU_TARGET_ARCHITECTURE} -nostdlib -static "--cuda-path=${LIBC_CUDA_ROOT}") @@ -647,14 +647,14 @@ function(add_libc_hermetic_test test_name) if(LIBC_TARGET_ARCHITECTURE_IS_AMDGPU) target_link_options(${fq_build_target_name} PRIVATE ${LIBC_COMPILE_OPTIONS_DEFAULT} - -mcpu=${LIBC_GPU_TARGET_ARCHITECTURE} -flto + -mcpu=${LIBC_GPU_TARGET_ARCHITECTURE} -flto -Wno-multi-gpu "-Wl,-mllvm,-amdgpu-lower-global-ctor-dtor=0" -nostdlib -static "-Wl,-mllvm,-amdhsa-code-object-version=${LIBC_GPU_CODE_OBJECT_VERSION}") elseif(LIBC_TARGET_ARCHITECTURE_IS_NVPTX) # We need to use the internal object versions for NVPTX. set(internal_suffix ".__internal__") target_link_options(${fq_build_target_name} PRIVATE - ${LIBC_COMPILE_OPTIONS_DEFAULT} + ${LIBC_COMPILE_OPTIONS_DEFAULT} -Wno-multi-gpu "-Wl,--suppress-stack-size-warning" -march=${LIBC_GPU_TARGET_ARCHITECTURE} -nostdlib -static "--cuda-path=${LIBC_CUDA_ROOT}") diff --git a/libc/include/llvm-libc-macros/math-macros.h b/libc/include/llvm-libc-macros/math-macros.h index 0a23647319f4..9f8edd954b7e 100644 --- a/libc/include/llvm-libc-macros/math-macros.h +++ b/libc/include/llvm-libc-macros/math-macros.h @@ -32,6 +32,8 @@ #define math_errhandling 0 #elif defined(__NO_MATH_ERRNO__) #define math_errhandling (MATH_ERREXCEPT) +#elif defined(__NVPTX__) || defined(__AMDGPU__) +#define math_errhandling (MATH_ERRNO) #else #define math_errhandling (MATH_ERRNO | MATH_ERREXCEPT) #endif diff --git a/libc/test/src/__support/CMakeLists.txt b/libc/test/src/__support/CMakeLists.txt index c5634866f839..7200ac276fe5 100644 --- a/libc/test/src/__support/CMakeLists.txt +++ b/libc/test/src/__support/CMakeLists.txt @@ -1,17 +1,14 @@ add_custom_target(libc-support-tests) -# FIXME: These tests are currently broken on the GPU. -if(NOT LIBC_TARGET_OS_IS_GPU) - add_libc_test( - blockstore_test - SUITE - libc-support-tests - SRCS - blockstore_test.cpp - DEPENDS - libc.src.__support.blockstore - ) -endif() +add_libc_test( + blockstore_test + SUITE + libc-support-tests + SRCS + blockstore_test.cpp + DEPENDS + libc.src.__support.blockstore +) add_libc_test( endian_test @@ -42,8 +39,6 @@ add_libc_test( DEPENDS libc.src.__support.high_precision_decimal libc.src.__support.uint128 - # FIXME Test segfaults on gfx90a GPU - UNIT_TEST_ONLY ) add_libc_test( diff --git a/libc/test/src/math/CMakeLists.txt b/libc/test/src/math/CMakeLists.txt index 81d2e1e55b55..ad7dfdb3dfd9 100644 --- a/libc/test/src/math/CMakeLists.txt +++ b/libc/test/src/math/CMakeLists.txt @@ -758,40 +758,37 @@ add_fp_unittest( libc.src.__support.FPUtil.basic_operations ) -# FIXME: These tests are currently broken for NVPTX. -if(NOT LIBC_TARGET_ARCHITECTURE_IS_NVPTX) - add_fp_unittest( - ilogb_test - SUITE - libc-math-unittests - SRCS - ilogb_test.cpp - HDRS - ILogbTest.h - DEPENDS - libc.include.math - libc.src.math.ilogb - libc.src.__support.CPP.limits - libc.src.__support.FPUtil.fp_bits - libc.src.__support.FPUtil.manipulation_functions - ) +add_fp_unittest( + ilogb_test + SUITE + libc-math-unittests + SRCS + ilogb_test.cpp + HDRS + ILogbTest.h + DEPENDS + libc.include.math + libc.src.math.ilogb + libc.src.__support.CPP.limits + libc.src.__support.FPUtil.fp_bits + libc.src.__support.FPUtil.manipulation_functions +) - add_fp_unittest( - ilogbf_test - SUITE - libc-math-unittests - SRCS - ilogbf_test.cpp - HDRS - ILogbTest.h - DEPENDS - libc.include.math - libc.src.math.ilogbf - libc.src.__support.CPP.limits - libc.src.__support.FPUtil.fp_bits - libc.src.__support.FPUtil.manipulation_functions - ) -endif() +add_fp_unittest( + ilogbf_test + SUITE + libc-math-unittests + SRCS + ilogbf_test.cpp + HDRS + ILogbTest.h + DEPENDS + libc.include.math + libc.src.math.ilogbf + libc.src.__support.CPP.limits + libc.src.__support.FPUtil.fp_bits + libc.src.__support.FPUtil.manipulation_functions +) add_fp_unittest( ilogbl_test @@ -989,92 +986,89 @@ add_fp_unittest( libc.src.__support.FPUtil.fp_bits ) -# FIXME: These tests are currently broken on the GPU. -if(NOT LIBC_TARGET_OS_IS_GPU) - add_fp_unittest( - fminf_test - SUITE - libc-math-unittests - SRCS - fminf_test.cpp - HDRS - FMinTest.h - DEPENDS - libc.include.math - libc.src.math.fminf - libc.src.__support.FPUtil.fp_bits - ) +add_fp_unittest( + fminf_test + SUITE + libc-math-unittests + SRCS + fminf_test.cpp + HDRS + FMinTest.h + DEPENDS + libc.include.math + libc.src.math.fminf + libc.src.__support.FPUtil.fp_bits +) - add_fp_unittest( - fmin_test - SUITE - libc-math-unittests - SRCS - fmin_test.cpp - HDRS - FMinTest.h - DEPENDS - libc.include.math - libc.src.math.fmin - libc.src.__support.FPUtil.fp_bits - ) +add_fp_unittest( + fmin_test + SUITE + libc-math-unittests + SRCS + fmin_test.cpp + HDRS + FMinTest.h + DEPENDS + libc.include.math + libc.src.math.fmin + libc.src.__support.FPUtil.fp_bits +) - add_fp_unittest( - fminl_test - SUITE - libc-math-unittests - SRCS - fminl_test.cpp - HDRS - FMinTest.h - DEPENDS - libc.include.math - libc.src.math.fminl - libc.src.__support.FPUtil.fp_bits - ) +add_fp_unittest( + fminl_test + SUITE + libc-math-unittests + SRCS + fminl_test.cpp + HDRS + FMinTest.h + DEPENDS + libc.include.math + libc.src.math.fminl + libc.src.__support.FPUtil.fp_bits +) - add_fp_unittest( - fmaxf_test - SUITE - libc-math-unittests - SRCS - fmaxf_test.cpp - HDRS - FMaxTest.h - DEPENDS - libc.include.math - libc.src.math.fmaxf - libc.src.__support.FPUtil.fp_bits - ) +add_fp_unittest( + fmaxf_test + SUITE + libc-math-unittests + SRCS + fmaxf_test.cpp + HDRS + FMaxTest.h + DEPENDS + libc.include.math + libc.src.math.fmaxf + libc.src.__support.FPUtil.fp_bits +) - add_fp_unittest( - fmax_test - SUITE - libc-math-unittests - SRCS - fmax_test.cpp - HDRS - FMaxTest.h - DEPENDS - libc.include.math - libc.src.math.fmax - libc.src.__support.FPUtil.fp_bits - ) +add_fp_unittest( + fmax_test + SUITE + libc-math-unittests + SRCS + fmax_test.cpp + HDRS + FMaxTest.h + DEPENDS + libc.include.math + libc.src.math.fmax + libc.src.__support.FPUtil.fp_bits +) - add_fp_unittest( - fmaxl_test - SUITE - libc-math-unittests - SRCS - fmaxl_test.cpp - HDRS - FMaxTest.h - DEPENDS - libc.include.math - libc.src.math.fmaxl - libc.src.__support.FPUtil.fp_bits - ) -endif() +add_fp_unittest( + fmaxl_test + SUITE + libc-math-unittests + SRCS + fmaxl_test.cpp + HDRS + FMaxTest.h + DEPENDS + libc.include.math + libc.src.math.fmaxl + libc.src.__support.FPUtil.fp_bits +) add_fp_unittest( sqrtf_test @@ -1234,38 +1228,35 @@ add_fp_unittest( libc.src.__support.FPUtil.fp_bits ) -# FIXME: These tests are currently spurious for NVPTX. -if(NOT LIBC_TARGET_ARCHITECTURE_IS_NVPTX) - add_fp_unittest( - nextafter_test - SUITE - libc-math-unittests - SRCS - nextafter_test.cpp - HDRS - NextAfterTest.h - DEPENDS - libc.include.math - libc.src.math.nextafter - libc.src.__support.FPUtil.basic_operations - libc.src.__support.FPUtil.fp_bits - ) +add_fp_unittest( + nextafter_test + SUITE + libc-math-unittests + SRCS + nextafter_test.cpp + HDRS + NextAfterTest.h + DEPENDS + libc.include.math + libc.src.math.nextafter + libc.src.__support.FPUtil.basic_operations + libc.src.__support.FPUtil.fp_bits +) - add_fp_unittest( - nextafterf_test - SUITE - libc-math-unittests - SRCS - nextafterf_test.cpp - HDRS - NextAfterTest.h - DEPENDS - libc.include.math - libc.src.math.nextafterf - libc.src.__support.FPUtil.basic_operations - libc.src.__support.FPUtil.fp_bits - ) -endif() +add_fp_unittest( + nextafterf_test + SUITE + libc-math-unittests + SRCS + nextafterf_test.cpp + HDRS + NextAfterTest.h + DEPENDS + libc.include.math + libc.src.math.nextafterf + libc.src.__support.FPUtil.basic_operations + libc.src.__support.FPUtil.fp_bits +) add_fp_unittest( nextafterl_test diff --git a/libc/test/src/math/smoke/CMakeLists.txt b/libc/test/src/math/smoke/CMakeLists.txt index 825000e1cb7a..be1810944495 100644 --- a/libc/test/src/math/smoke/CMakeLists.txt +++ b/libc/test/src/math/smoke/CMakeLists.txt @@ -758,38 +758,35 @@ add_fp_unittest( libc.src.math.frexpf128 ) -# FIXME: These tests are currently broken for NVPTX. -if(NOT LIBC_TARGET_ARCHITECTURE_IS_NVPTX) - add_fp_unittest( - ilogb_test - SUITE - libc-math-smoke-tests - SRCS - ilogb_test.cpp - HDRS - ILogbTest.h - DEPENDS - libc.src.math.ilogb - libc.src.__support.CPP.limits - libc.src.__support.FPUtil.fp_bits - libc.src.__support.FPUtil.manipulation_functions - ) +add_fp_unittest( + ilogb_test + SUITE + libc-math-smoke-tests + SRCS + ilogb_test.cpp + HDRS + ILogbTest.h + DEPENDS + libc.src.math.ilogb + libc.src.__support.CPP.limits + libc.src.__support.FPUtil.fp_bits + libc.src.__support.FPUtil.manipulation_functions +) - add_fp_unittest( - ilogbf_test - SUITE - libc-math-smoke-tests - SRCS - ilogbf_test.cpp - HDRS - ILogbTest.h - DEPENDS - libc.src.math.ilogbf - libc.src.__support.CPP.limits - libc.src.__support.FPUtil.fp_bits - libc.src.__support.FPUtil.manipulation_functions - ) -endif() +add_fp_unittest( + ilogbf_test + SUITE + libc-math-smoke-tests + SRCS + ilogbf_test.cpp + HDRS + ILogbTest.h + DEPENDS + libc.src.math.ilogbf + libc.src.__support.CPP.limits + libc.src.__support.FPUtil.fp_bits + libc.src.__support.FPUtil.manipulation_functions +) add_fp_unittest( ilogbl_test @@ -1417,38 +1414,35 @@ add_fp_unittest( UNIT_TEST_ONLY ) -# FIXME: These tests are currently spurious for NVPTX. -if(NOT LIBC_TARGET_ARCHITECTURE_IS_NVPTX) - add_fp_unittest( - nextafter_test - SUITE - libc-math-smoke-tests - SRCS - nextafter_test.cpp - HDRS - NextAfterTest.h - DEPENDS - libc.include.math - libc.src.math.nextafter - libc.src.__support.FPUtil.basic_operations - libc.src.__support.FPUtil.fp_bits - ) +add_fp_unittest( + nextafter_test + SUITE + libc-math-smoke-tests + SRCS + nextafter_test.cpp + HDRS + NextAfterTest.h + DEPENDS + libc.include.math + libc.src.math.nextafter + libc.src.__support.FPUtil.basic_operations + libc.src.__support.FPUtil.fp_bits +) - add_fp_unittest( - nextafterf_test - SUITE - libc-math-smoke-tests - SRCS - nextafterf_test.cpp - HDRS - NextAfterTest.h - DEPENDS - libc.include.math - libc.src.math.nextafterf - libc.src.__support.FPUtil.basic_operations - libc.src.__support.FPUtil.fp_bits - ) -endif() +add_fp_unittest( + nextafterf_test + SUITE + libc-math-smoke-tests + SRCS + nextafterf_test.cpp + HDRS + NextAfterTest.h + DEPENDS + libc.include.math + libc.src.math.nextafterf + libc.src.__support.FPUtil.basic_operations + libc.src.__support.FPUtil.fp_bits +) add_fp_unittest( nextafterl_test diff --git a/libc/test/src/stdlib/CMakeLists.txt b/libc/test/src/stdlib/CMakeLists.txt index 5826cfe8d4ca..5488a61c4ef1 100644 --- a/libc/test/src/stdlib/CMakeLists.txt +++ b/libc/test/src/stdlib/CMakeLists.txt @@ -54,20 +54,17 @@ add_libc_test( libc.src.stdlib.atoll ) -# This fails on NVPTX where the output value is one-off of the expected value. -if(NOT LIBC_TARGET_ARCHITECTURE_IS_NVPTX) - add_fp_unittest( - strtod_test - SUITE - libc-stdlib-tests - SRCS - strtod_test.cpp - DEPENDS - libc.src.errno.errno - libc.src.stdlib.strtod - libc.src.__support.FPUtil.fenv_impl - ) -endif() +add_fp_unittest( + strtod_test + SUITE + libc-stdlib-tests + SRCS + strtod_test.cpp + DEPENDS + libc.src.errno.errno + libc.src.stdlib.strtod + libc.src.__support.FPUtil.fenv_impl +) add_fp_unittest( strtof_test @@ -126,20 +123,17 @@ add_libc_test( .strtol_test_support ) -# This fails on NVPTX where the output value is one-off of the expected value. -if(NOT LIBC_TARGET_ARCHITECTURE_IS_NVPTX) - add_libc_test( - strtold_test - SUITE - libc-stdlib-tests - SRCS - strtold_test.cpp - DEPENDS - libc.src.errno.errno - libc.src.__support.uint128 - libc.src.stdlib.strtold - ) -endif() +add_libc_test( + strtold_test + SUITE + libc-stdlib-tests + SRCS + strtold_test.cpp + DEPENDS + libc.src.errno.errno + libc.src.__support.uint128 + libc.src.stdlib.strtold +) add_libc_test( strtoll_test diff --git a/libc/test/src/time/CMakeLists.txt b/libc/test/src/time/CMakeLists.txt index ebb0998feb23..51cacef0a62f 100644 --- a/libc/test/src/time/CMakeLists.txt +++ b/libc/test/src/time/CMakeLists.txt @@ -102,21 +102,17 @@ add_libc_unittest( libc.src.__support.CPP.limits ) -# Sleeping is not supported on older NVPTX architectures. -set(unsupported_architectures "sm_35;sm_37;sm_50;sm_52;sm_53;sm_60;sm_61;sm_62") -if (NOT ("${LIBC_GPU_TARGET_ARCHITECTURE}" IN_LIST unsupported_architectures)) - add_libc_test( - nanosleep_test - SUITE - libc_time_unittests - SRCS - nanosleep_test.cpp - DEPENDS - libc.include.time - libc.src.time.nanosleep - libc.src.errno.errno - ) -endif() +add_libc_test( + nanosleep_test + SUITE + libc_time_unittests + SRCS + nanosleep_test.cpp + DEPENDS + libc.include.time + libc.src.time.nanosleep + libc.src.errno.errno +) add_libc_unittest( time_test diff --git a/lld/MachO/Driver.cpp b/lld/MachO/Driver.cpp index a57f60c5eed3..018ceec97f20 100644 --- a/lld/MachO/Driver.cpp +++ b/lld/MachO/Driver.cpp @@ -691,6 +691,8 @@ static PlatformVersion parsePlatformVersion(const Arg *arg) { .Cases("tvos-simulator", "8", PLATFORM_TVOSSIMULATOR) .Cases("watchos-simulator", "9", PLATFORM_WATCHOSSIMULATOR) .Cases("driverkit", "10", PLATFORM_DRIVERKIT) + .Cases("xros", "11", PLATFORM_XROS) + .Cases("xros-simulator", "12", PLATFORM_XROS_SIMULATOR) .Default(PLATFORM_UNKNOWN); if (platformVersion.platform == PLATFORM_UNKNOWN) error(Twine("malformed platform: ") + platformStr); @@ -985,6 +987,8 @@ PlatformType macho::removeSimulator(PlatformType platform) { return PLATFORM_TVOS; case PLATFORM_WATCHOSSIMULATOR: return PLATFORM_WATCHOS; + case PLATFORM_XROS_SIMULATOR: + return PLATFORM_XROS; default: return platform; } @@ -1001,15 +1005,17 @@ static bool shouldAdhocSignByDefault(Architecture arch, PlatformType platform) { return platform == PLATFORM_MACOS || platform == PLATFORM_IOSSIMULATOR || platform == PLATFORM_TVOSSIMULATOR || - platform == PLATFORM_WATCHOSSIMULATOR; + platform == PLATFORM_WATCHOSSIMULATOR || + platform == PLATFORM_XROS_SIMULATOR; } static bool dataConstDefault(const InputArgList &args) { - static const std::array<std::pair<PlatformType, VersionTuple>, 5> minVersion = + static const std::array<std::pair<PlatformType, VersionTuple>, 6> minVersion = {{{PLATFORM_MACOS, VersionTuple(10, 15)}, {PLATFORM_IOS, VersionTuple(13, 0)}, {PLATFORM_TVOS, VersionTuple(13, 0)}, {PLATFORM_WATCHOS, VersionTuple(6, 0)}, + {PLATFORM_XROS, VersionTuple(1, 0)}, {PLATFORM_BRIDGEOS, VersionTuple(4, 0)}}}; PlatformType platform = removeSimulator(config->platformInfo.target.Platform); auto it = llvm::find_if(minVersion, @@ -1045,11 +1051,12 @@ static bool shouldEmitChainedFixups(const InputArgList &args) { bool isRequested = arg != nullptr; // Version numbers taken from the Xcode 13.3 release notes. - static const std::array<std::pair<PlatformType, VersionTuple>, 4> minVersion = + static const std::array<std::pair<PlatformType, VersionTuple>, 5> minVersion = {{{PLATFORM_MACOS, VersionTuple(11, 0)}, {PLATFORM_IOS, VersionTuple(13, 4)}, {PLATFORM_TVOS, VersionTuple(14, 0)}, - {PLATFORM_WATCHOS, VersionTuple(7, 0)}}}; + {PLATFORM_WATCHOS, VersionTuple(7, 0)}, + {PLATFORM_XROS, VersionTuple(1, 0)}}}; PlatformType platform = removeSimulator(config->platformInfo.target.Platform); auto it = llvm::find_if(minVersion, [&](const auto &p) { return p.first == platform; }); @@ -1688,8 +1695,8 @@ bool link(ArrayRef<const char *> argsArr, llvm::raw_ostream &stdoutOS, if (args.getLastArg(OPT_reproducible)) config->zeroModTime = true; - std::array<PlatformType, 3> encryptablePlatforms{ - PLATFORM_IOS, PLATFORM_WATCHOS, PLATFORM_TVOS}; + std::array<PlatformType, 4> encryptablePlatforms{ + PLATFORM_IOS, PLATFORM_WATCHOS, PLATFORM_TVOS, PLATFORM_XROS}; config->emitEncryptionInfo = args.hasFlag(OPT_encryptable, OPT_no_encryption, is_contained(encryptablePlatforms, config->platform())); diff --git a/lld/MachO/Options.td b/lld/MachO/Options.td index 01e73b789f9a..a524e4a4c508 100644 --- a/lld/MachO/Options.td +++ b/lld/MachO/Options.td @@ -377,7 +377,7 @@ def grp_version : OptionGroup<"version">, HelpText<"VERSION TARGETING">; def platform_version : MultiArg<["-"], "platform_version", 3>, MetaVarName<"<platform> <min_version> <sdk_version>">, - HelpText<"Platform (e.g., macos, ios, tvos, watchos, bridgeos, mac-catalyst, ios-sim, tvos-sim, watchos-sim, driverkit) and version numbers">, + HelpText<"Platform (e.g., macos, ios, tvos, watchos, xros, bridgeos, mac-catalyst, ios-sim, tvos-sim, watchos-sim, xros-sim, driverkit) and version numbers">, Group<grp_version>; def sdk_version : Separate<["-"], "sdk_version">, HelpText<"This option is undocumented in ld64">, diff --git a/lld/test/MachO/lc-build-version.s b/lld/test/MachO/lc-build-version.s index 7b78f803428a..1fd7078919b1 100644 --- a/lld/test/MachO/lc-build-version.s +++ b/lld/test/MachO/lc-build-version.s @@ -64,6 +64,13 @@ # WATCHOS-4-0: cmd LC_VERSION_MIN_WATCHOS +# RUN: %no-arg-lld -arch x86_64 -platform_version xros 1.0 1.1 -o %t.xros-1-0 %t.o +# RUN: llvm-objdump --macho --all-headers %t.xros-1-0 | FileCheck %s --check-prefix=XROS-1-0 +# RUN: %no-arg-lld -arch x86_64 -platform_version xros-simulator 1.0 1.1 -o %t.xros-sim-1-0 %t.o +# RUN: llvm-objdump --macho --all-headers %t.xros-sim-1-0 | FileCheck %s --check-prefix=XROS-1-0 + +# XROS-1-0: cmd LC_BUILD_VERSION + .text .global _main _main: diff --git a/lld/test/MachO/platform-version.s b/lld/test/MachO/platform-version.s index 047aea02fcde..57fbae62b2ff 100644 --- a/lld/test/MachO/platform-version.s +++ b/lld/test/MachO/platform-version.s @@ -55,7 +55,7 @@ # RUN: -platform_version 0 1 5 \ # RUN: | FileCheck --check-prefix=FAIL-PLATFORM %s # RUN: not %no-arg-lld -arch x86_64 -o %t %t.o 2>&1 \ -# RUN: -platform_version 11 1 5 \ +# RUN: -platform_version 13 1 5 \ # RUN: | FileCheck --check-prefix=FAIL-PLATFORM %s # FAIL-PLATFORM: malformed platform: {{.*}} # FAIL-PLATFORM-NOT: malformed {{minimum|sdk}} version: {{.*}} diff --git a/llvm/docs/CommandGuide/llvm-exegesis.rst b/llvm/docs/CommandGuide/llvm-exegesis.rst index 9e3c19078f1c..fdf17c7fe412 100644 --- a/llvm/docs/CommandGuide/llvm-exegesis.rst +++ b/llvm/docs/CommandGuide/llvm-exegesis.rst @@ -89,6 +89,14 @@ properly. annotation requires the subprocess execution mode. This is useful in cases where the memory accessed by the snippet depends on the location of the snippet, like RIP-relative addressing. +* `LLVM-EXEGESIS-LOOP-REGISTER <register name>` - This annotation specifies + the loop register to use for keeping track of the current iteration when + using the loop repetition mode. :program:`llvm-exegesis` needs to keep track + of the current loop iteration within the loop repetition mode in a performant + manner (i.e., no memory accesses), and uses a register to do this. This register + has an architecture specific default (e.g., `R8` on X86), but this might conflict + with some snippets. This annotation allows changing the register to prevent + interference between the loop index register and the snippet. EXAMPLE 1: benchmarking instructions ------------------------------------ diff --git a/llvm/docs/CommandGuide/llvm-objdump.rst b/llvm/docs/CommandGuide/llvm-objdump.rst index 959452a74b23..7f8def756c69 100644 --- a/llvm/docs/CommandGuide/llvm-objdump.rst +++ b/llvm/docs/CommandGuide/llvm-objdump.rst @@ -271,7 +271,12 @@ OPTIONS When printing a PC-relative global symbol reference, print it as an offset from the leading symbol. - When a bb-address-map section is present (i.e., the object file is built with ``-fbasic-block-sections=labels``), labels are retrieved from that section instead. + When a bb-address-map section is present (i.e., the object file is built with + ``-fbasic-block-sections=labels``), labels are retrieved from that section + instead. If a pgo-analysis-map is present alongside the bb-address-map, any + available analyses are printed after the relevant block label. By default, + any analysis with a special representation (i.e. BlockFrequency, + BranchProbability, etc) are printed as raw hex values. Only works with PowerPC objects or X86 linked images. @@ -291,6 +296,15 @@ OPTIONS cmp eax, dword ptr <g> jge <L0> +.. option:: --pretty-pgo-analysis-map + + When using :option:`--symbolize-operands` with bb-address-map and + pgo-analysis-map, print analyses using the same format as their analysis + passes would. An example of pretty format would be printing block frequencies + relative to the entry block, the same as BFI. + + Only works when :option:`--symbolize-operands` is enabled. + .. option:: --triple=<string> Target triple to disassemble for, see ``--version`` for available targets. diff --git a/llvm/docs/CommandGuide/llvm-readobj.rst b/llvm/docs/CommandGuide/llvm-readobj.rst index 6d78a0387234..09dabb28cfa7 100644 --- a/llvm/docs/CommandGuide/llvm-readobj.rst +++ b/llvm/docs/CommandGuide/llvm-readobj.rst @@ -164,6 +164,17 @@ The following options are implemented only for the ELF file format. Display the contents of the basic block address map section(s), which contain the address of each function, along with the relative offset of each basic block. + When pgo analysis maps are present, all analyses are printed as their raw + value. + +.. option:: --pretty-pgo-analysis-map + + When pgo analysis maps are present in the basic block address map section(s), + analyses with special formats (i.e. BlockFrequency, BranchProbability, etc) + are printed using the same format as their respective analysis pass. + + Requires :option:`--bb-addr-map` to have an effect. + .. option:: --demangle, -C Display demangled symbol names in the output. diff --git a/llvm/include/llvm/Analysis/BlockFrequencyInfoImpl.h b/llvm/include/llvm/Analysis/BlockFrequencyInfoImpl.h index 8acb75e87254..4aa922635c37 100644 --- a/llvm/include/llvm/Analysis/BlockFrequencyInfoImpl.h +++ b/llvm/include/llvm/Analysis/BlockFrequencyInfoImpl.h @@ -539,9 +539,6 @@ public: } }; -void printBlockFreqImpl(raw_ostream &OS, BlockFrequency EntryFreq, - BlockFrequency Freq); - namespace bfi_detail { template <class BlockT> struct TypeMap {}; diff --git a/llvm/include/llvm/Support/BlockFrequency.h b/llvm/include/llvm/Support/BlockFrequency.h index 8b172ee486aa..aeab99615a95 100644 --- a/llvm/include/llvm/Support/BlockFrequency.h +++ b/llvm/include/llvm/Support/BlockFrequency.h @@ -19,6 +19,7 @@ namespace llvm { +class raw_ostream; class BranchProbability; // This class represents Block Frequency as a 64-bit value. @@ -119,6 +120,9 @@ public: } }; +void printRelativeBlockFreq(raw_ostream &OS, BlockFrequency EntryFreq, + BlockFrequency Freq); + } // namespace llvm #endif diff --git a/llvm/lib/Analysis/BlockFrequencyInfo.cpp b/llvm/lib/Analysis/BlockFrequencyInfo.cpp index 96c9bfa0e372..ebad8388cbe4 100644 --- a/llvm/lib/Analysis/BlockFrequencyInfo.cpp +++ b/llvm/lib/Analysis/BlockFrequencyInfo.cpp @@ -284,7 +284,7 @@ void BlockFrequencyInfo::verifyMatch(BlockFrequencyInfo &Other) const { Printable llvm::printBlockFreq(const BlockFrequencyInfo &BFI, BlockFrequency Freq) { return Printable([&BFI, Freq](raw_ostream &OS) { - printBlockFreqImpl(OS, BFI.getEntryFreq(), Freq); + printRelativeBlockFreq(OS, BFI.getEntryFreq(), Freq); }); } diff --git a/llvm/lib/Analysis/BlockFrequencyInfoImpl.cpp b/llvm/lib/Analysis/BlockFrequencyInfoImpl.cpp index ae08d56ef098..9f6e53ba15b6 100644 --- a/llvm/lib/Analysis/BlockFrequencyInfoImpl.cpp +++ b/llvm/lib/Analysis/BlockFrequencyInfoImpl.cpp @@ -634,21 +634,6 @@ BlockFrequencyInfoImplBase::getLoopName(const LoopData &Loop) const { return getBlockName(Loop.getHeader()) + (Loop.isIrreducible() ? "**" : "*"); } -void llvm::printBlockFreqImpl(raw_ostream &OS, BlockFrequency EntryFreq, - BlockFrequency Freq) { - if (Freq == BlockFrequency(0)) { - OS << "0"; - return; - } - if (EntryFreq == BlockFrequency(0)) { - OS << "<invalid BFI>"; - return; - } - Scaled64 Block(Freq.getFrequency(), 0); - Scaled64 Entry(EntryFreq.getFrequency(), 0); - OS << Block / Entry; -} - void IrreducibleGraph::addNodesInLoop(const BFIBase::LoopData &OuterLoop) { Start = OuterLoop.getHeader(); Nodes.reserve(OuterLoop.Nodes.size()); diff --git a/llvm/lib/CodeGen/MachineBlockFrequencyInfo.cpp b/llvm/lib/CodeGen/MachineBlockFrequencyInfo.cpp index 7ee72e214426..cbebdd87398e 100644 --- a/llvm/lib/CodeGen/MachineBlockFrequencyInfo.cpp +++ b/llvm/lib/CodeGen/MachineBlockFrequencyInfo.cpp @@ -280,7 +280,7 @@ BlockFrequency MachineBlockFrequencyInfo::getEntryFreq() const { Printable llvm::printBlockFreq(const MachineBlockFrequencyInfo &MBFI, BlockFrequency Freq) { return Printable([&MBFI, Freq](raw_ostream &OS) { - printBlockFreqImpl(OS, MBFI.getEntryFreq(), Freq); + printRelativeBlockFreq(OS, MBFI.getEntryFreq(), Freq); }); } diff --git a/llvm/lib/Support/BlockFrequency.cpp b/llvm/lib/Support/BlockFrequency.cpp index 329f1e12cdc2..7d5498e7cb99 100644 --- a/llvm/lib/Support/BlockFrequency.cpp +++ b/llvm/lib/Support/BlockFrequency.cpp @@ -13,6 +13,8 @@ #include "llvm/Support/BlockFrequency.h" #include "llvm/Support/BranchProbability.h" #include "llvm/Support/MathExtras.h" +#include "llvm/Support/ScaledNumber.h" +#include "llvm/Support/raw_ostream.h" using namespace llvm; @@ -45,3 +47,18 @@ std::optional<BlockFrequency> BlockFrequency::mul(uint64_t Factor) const { return {}; return BlockFrequency(ResultFrequency); } + +void llvm::printRelativeBlockFreq(raw_ostream &OS, BlockFrequency EntryFreq, + BlockFrequency Freq) { + if (Freq == BlockFrequency(0)) { + OS << "0"; + return; + } + if (EntryFreq == BlockFrequency(0)) { + OS << "<invalid BFI>"; + return; + } + ScaledNumber<uint64_t> Block(Freq.getFrequency(), 0); + ScaledNumber<uint64_t> Entry(EntryFreq.getFrequency(), 0); + OS << Block / Entry; +} diff --git a/llvm/lib/Target/AArch64/AArch64InstrInfo.td b/llvm/lib/Target/AArch64/AArch64InstrInfo.td index e73bc0d89e4c..b01a8cd00025 100644 --- a/llvm/lib/Target/AArch64/AArch64InstrInfo.td +++ b/llvm/lib/Target/AArch64/AArch64InstrInfo.td @@ -6667,31 +6667,29 @@ def : Pat<(vector_extract (v8bf16 V128:$Rn), VectorIndexH:$idx), // All concat_vectors operations are canonicalised to act on i64 vectors for // AArch64. In the general case we need an instruction, which had just as well be // INS. -class ConcatPat<ValueType DstTy, ValueType SrcTy> - : Pat<(DstTy (concat_vectors (SrcTy V64:$Rd), V64:$Rn)), - (INSvi64lane (INSERT_SUBREG (IMPLICIT_DEF), V64:$Rd, dsub), 1, - (INSERT_SUBREG (IMPLICIT_DEF), V64:$Rn, dsub), 0)>; - -def : ConcatPat<v2i64, v1i64>; -def : ConcatPat<v2f64, v1f64>; -def : ConcatPat<v4i32, v2i32>; -def : ConcatPat<v4f32, v2f32>; -def : ConcatPat<v8i16, v4i16>; -def : ConcatPat<v8f16, v4f16>; -def : ConcatPat<v8bf16, v4bf16>; -def : ConcatPat<v16i8, v8i8>; - -// If the high lanes are undef, though, we can just ignore them: -class ConcatUndefPat<ValueType DstTy, ValueType SrcTy> - : Pat<(DstTy (concat_vectors (SrcTy V64:$Rn), undef)), - (INSERT_SUBREG (IMPLICIT_DEF), V64:$Rn, dsub)>; - -def : ConcatUndefPat<v2i64, v1i64>; -def : ConcatUndefPat<v2f64, v1f64>; -def : ConcatUndefPat<v4i32, v2i32>; -def : ConcatUndefPat<v4f32, v2f32>; -def : ConcatUndefPat<v8i16, v4i16>; -def : ConcatUndefPat<v16i8, v8i8>; +multiclass ConcatPat<ValueType DstTy, ValueType SrcTy> { + def : Pat<(DstTy (concat_vectors (SrcTy V64:$Rd), V64:$Rn)), + (INSvi64lane (INSERT_SUBREG (IMPLICIT_DEF), V64:$Rd, dsub), 1, + (INSERT_SUBREG (IMPLICIT_DEF), V64:$Rn, dsub), 0)>; + + // If the high lanes are zero we can instead emit a d->d register mov, which + // will implicitly clear the upper bits. + def : Pat<(DstTy (concat_vectors (SrcTy V64:$Rn), immAllZerosV)), + (SUBREG_TO_REG (i64 0), (FMOVDr V64:$Rn), dsub)>; + + // If the high lanes are undef we can just ignore them: + def : Pat<(DstTy (concat_vectors (SrcTy V64:$Rn), undef)), + (INSERT_SUBREG (IMPLICIT_DEF), V64:$Rn, dsub)>; +} + +defm : ConcatPat<v2i64, v1i64>; +defm : ConcatPat<v2f64, v1f64>; +defm : ConcatPat<v4i32, v2i32>; +defm : ConcatPat<v4f32, v2f32>; +defm : ConcatPat<v8i16, v4i16>; +defm : ConcatPat<v8f16, v4f16>; +defm : ConcatPat<v8bf16, v4bf16>; +defm : ConcatPat<v16i8, v8i8>; //---------------------------------------------------------------------------- // AdvSIMD across lanes instructions diff --git a/llvm/lib/Target/AArch64/AArch64MIPeepholeOpt.cpp b/llvm/lib/Target/AArch64/AArch64MIPeepholeOpt.cpp index 87aa3b98d938..6865850cf04f 100644 --- a/llvm/lib/Target/AArch64/AArch64MIPeepholeOpt.cpp +++ b/llvm/lib/Target/AArch64/AArch64MIPeepholeOpt.cpp @@ -127,6 +127,7 @@ struct AArch64MIPeepholeOpt : public MachineFunctionPass { bool visitINSERT(MachineInstr &MI); bool visitINSviGPR(MachineInstr &MI, unsigned Opc); bool visitINSvi64lane(MachineInstr &MI); + bool visitFMOVDr(MachineInstr &MI); bool runOnMachineFunction(MachineFunction &MF) override; StringRef getPassName() const override { @@ -670,6 +671,23 @@ bool AArch64MIPeepholeOpt::visitINSvi64lane(MachineInstr &MI) { return true; } +bool AArch64MIPeepholeOpt::visitFMOVDr(MachineInstr &MI) { + // An FMOVDr sets the high 64-bits to zero implicitly, similar to ORR for GPR. + MachineInstr *Low64MI = MRI->getUniqueVRegDef(MI.getOperand(1).getReg()); + if (!Low64MI || !is64bitDefwithZeroHigh64bit(Low64MI, MRI)) + return false; + + // Let's remove MIs for high 64-bits. + Register OldDef = MI.getOperand(0).getReg(); + Register NewDef = MI.getOperand(1).getReg(); + MRI->constrainRegClass(NewDef, MRI->getRegClass(OldDef)); + MRI->replaceRegWith(OldDef, NewDef); + LLVM_DEBUG(dbgs() << "Removed: " << MI << "\n"); + MI.eraseFromParent(); + + return true; +} + bool AArch64MIPeepholeOpt::runOnMachineFunction(MachineFunction &MF) { if (skipFunction(MF.getFunction())) return false; @@ -748,6 +766,9 @@ bool AArch64MIPeepholeOpt::runOnMachineFunction(MachineFunction &MF) { case AArch64::INSvi64lane: Changed |= visitINSvi64lane(MI); break; + case AArch64::FMOVDr: + Changed |= visitFMOVDr(MI); + break; } } } diff --git a/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp b/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp index 8a2864a07873..5d5c4ea57ed5 100644 --- a/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp +++ b/llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp @@ -1957,6 +1957,10 @@ bool ModuleAddressSanitizer::shouldInstrumentGlobal(GlobalVariable *G) const { // On COFF, don't instrument non-ODR linkages. if (G->isInterposable()) return false; + // If the global has AvailableExternally linkage, then it is not in this + // module, which means it does not need to be instrumented. + if (G->hasAvailableExternallyLinkage()) + return false; } // If a comdat is present, it must have a selection kind that implies ODR diff --git a/llvm/test/CodeGen/AArch64/implicitly-set-zero-high-64-bits.ll b/llvm/test/CodeGen/AArch64/implicitly-set-zero-high-64-bits.ll index 1eb9eab1c21e..a949eaac5cfa 100644 --- a/llvm/test/CodeGen/AArch64/implicitly-set-zero-high-64-bits.ll +++ b/llvm/test/CodeGen/AArch64/implicitly-set-zero-high-64-bits.ll @@ -137,9 +137,7 @@ entry: define <16 x i8> @insertzero_v8i8(<8 x i8> %a) { ; CHECK-LABEL: insertzero_v8i8: ; CHECK: // %bb.0: // %entry -; CHECK-NEXT: movi v1.2d, #0000000000000000 -; CHECK-NEXT: // kill: def $d0 killed $d0 def $q0 -; CHECK-NEXT: mov v0.d[1], v1.d[0] +; CHECK-NEXT: fmov d0, d0 ; CHECK-NEXT: ret entry: %shuffle.i = shufflevector <8 x i8> %a, <8 x i8> zeroinitializer, <16 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7, i32 8, i32 9, i32 10, i32 11, i32 12, i32 13, i32 14, i32 15> @@ -149,9 +147,7 @@ entry: define <8 x i16> @insertzero_v4i16(<4 x i16> %a) { ; CHECK-LABEL: insertzero_v4i16: ; CHECK: // %bb.0: // %entry -; CHECK-NEXT: movi v1.2d, #0000000000000000 -; CHECK-NEXT: // kill: def $d0 killed $d0 def $q0 -; CHECK-NEXT: mov v0.d[1], v1.d[0] +; CHECK-NEXT: fmov d0, d0 ; CHECK-NEXT: ret entry: %shuffle.i = shufflevector <4 x i16> %a, <4 x i16> zeroinitializer, <8 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7> @@ -161,9 +157,7 @@ entry: define <4 x i32> @insertzero_v2i32(<2 x i32> %a) { ; CHECK-LABEL: insertzero_v2i32: ; CHECK: // %bb.0: // %entry -; CHECK-NEXT: movi v1.2d, #0000000000000000 -; CHECK-NEXT: // kill: def $d0 killed $d0 def $q0 -; CHECK-NEXT: mov v0.d[1], v1.d[0] +; CHECK-NEXT: fmov d0, d0 ; CHECK-NEXT: ret entry: %shuffle.i = shufflevector <2 x i32> %a, <2 x i32> zeroinitializer, <4 x i32> <i32 0, i32 1, i32 2, i32 3> @@ -173,9 +167,7 @@ entry: define <2 x i64> @insertzero_v1i64(<1 x i64> %a) { ; CHECK-LABEL: insertzero_v1i64: ; CHECK: // %bb.0: // %entry -; CHECK-NEXT: movi v1.2d, #0000000000000000 -; CHECK-NEXT: // kill: def $d0 killed $d0 def $q0 -; CHECK-NEXT: mov v0.d[1], v1.d[0] +; CHECK-NEXT: fmov d0, d0 ; CHECK-NEXT: ret entry: %shuffle.i = shufflevector <1 x i64> %a, <1 x i64> zeroinitializer, <2 x i32> <i32 0, i32 1> @@ -185,9 +177,7 @@ entry: define <8 x half> @insertzero_v4f16(<4 x half> %a) { ; CHECK-LABEL: insertzero_v4f16: ; CHECK: // %bb.0: // %entry -; CHECK-NEXT: movi d1, #0000000000000000 -; CHECK-NEXT: // kill: def $d0 killed $d0 def $q0 -; CHECK-NEXT: mov v0.d[1], v1.d[0] +; CHECK-NEXT: fmov d0, d0 ; CHECK-NEXT: ret entry: %shuffle.i = shufflevector <4 x half> %a, <4 x half> zeroinitializer, <8 x i32> <i32 0, i32 1, i32 2, i32 3, i32 4, i32 5, i32 6, i32 7> @@ -210,9 +200,7 @@ entry: define <4 x float> @insertzero_v2f32(<2 x float> %a) { ; CHECK-LABEL: insertzero_v2f32: ; CHECK: // %bb.0: // %entry -; CHECK-NEXT: movi d1, #0000000000000000 -; CHECK-NEXT: // kill: def $d0 killed $d0 def $q0 -; CHECK-NEXT: mov v0.d[1], v1.d[0] +; CHECK-NEXT: fmov d0, d0 ; CHECK-NEXT: ret entry: %shuffle.i = shufflevector <2 x float> %a, <2 x float> zeroinitializer, <4 x i32> <i32 0, i32 1, i32 2, i32 3> @@ -222,9 +210,7 @@ entry: define <2 x double> @insertzero_v1f64(<1 x double> %a) { ; CHECK-LABEL: insertzero_v1f64: ; CHECK: // %bb.0: // %entry -; CHECK-NEXT: movi d1, #0000000000000000 -; CHECK-NEXT: // kill: def $d0 killed $d0 def $q0 -; CHECK-NEXT: mov v0.d[1], v1.d[0] +; CHECK-NEXT: fmov d0, d0 ; CHECK-NEXT: ret entry: %shuffle.i = shufflevector <1 x double> %a, <1 x double> zeroinitializer, <2 x i32> <i32 0, i32 1> diff --git a/llvm/test/CodeGen/AArch64/vecreduce-add.ll b/llvm/test/CodeGen/AArch64/vecreduce-add.ll index 86dd1bdd511e..66b49466cc73 100644 --- a/llvm/test/CodeGen/AArch64/vecreduce-add.ll +++ b/llvm/test/CodeGen/AArch64/vecreduce-add.ll @@ -2182,8 +2182,8 @@ define i32 @test_udot_v24i8(ptr %p1, ptr %p2) { ; CHECK-GI-DOT-NEXT: ldr b5, [x0, #15] ; CHECK-GI-DOT-NEXT: mov v2.b[14], v6.b[0] ; CHECK-GI-DOT-NEXT: ldr b6, [x1, #15] -; CHECK-GI-DOT-NEXT: mov v3.d[1], v0.d[0] -; CHECK-GI-DOT-NEXT: mov v4.d[1], v0.d[0] +; CHECK-GI-DOT-NEXT: fmov d3, d3 +; CHECK-GI-DOT-NEXT: fmov d4, d4 ; CHECK-GI-DOT-NEXT: mov v1.b[15], v5.b[0] ; CHECK-GI-DOT-NEXT: movi v5.2d, #0000000000000000 ; CHECK-GI-DOT-NEXT: mov v2.b[15], v6.b[0] @@ -2760,8 +2760,8 @@ define i32 @test_sdot_v24i8(ptr %p1, ptr %p2) { ; CHECK-GI-DOT-NEXT: ldr b5, [x0, #15] ; CHECK-GI-DOT-NEXT: mov v2.b[14], v6.b[0] ; CHECK-GI-DOT-NEXT: ldr b6, [x1, #15] -; CHECK-GI-DOT-NEXT: mov v3.d[1], v0.d[0] -; CHECK-GI-DOT-NEXT: mov v4.d[1], v0.d[0] +; CHECK-GI-DOT-NEXT: fmov d3, d3 +; CHECK-GI-DOT-NEXT: fmov d4, d4 ; CHECK-GI-DOT-NEXT: mov v1.b[15], v5.b[0] ; CHECK-GI-DOT-NEXT: movi v5.2d, #0000000000000000 ; CHECK-GI-DOT-NEXT: mov v2.b[15], v6.b[0] diff --git a/llvm/test/Instrumentation/AddressSanitizer/do-not-instrument-globals-windows.ll b/llvm/test/Instrumentation/AddressSanitizer/do-not-instrument-globals-windows.ll new file mode 100644 index 000000000000..c143f69f126a --- /dev/null +++ b/llvm/test/Instrumentation/AddressSanitizer/do-not-instrument-globals-windows.ll @@ -0,0 +1,10 @@ +; This test checks that we are not instrumenting unnecessary globals +; RUN: opt < %s -passes=asan -S | FileCheck %s + +target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128" +target triple = "x86_64-pc-windows-msvc" + +@v_available_externally = available_externally global i32 zeroinitializer +; CHECK-NOT: {{asan_gen.*v_available_externally}} + +; CHECK: @asan.module_ctor diff --git a/llvm/test/tools/llvm-exegesis/X86/latency/loop-register.s b/llvm/test/tools/llvm-exegesis/X86/latency/loop-register.s new file mode 100644 index 000000000000..81ca75251381 --- /dev/null +++ b/llvm/test/tools/llvm-exegesis/X86/latency/loop-register.s @@ -0,0 +1,12 @@ +# REQUIRES: exegesis-can-measure-latency, x86_64-linux + +# Test that specifying the loop register to use works as expected. + +# RUN: llvm-exegesis -mtriple=x86_64-unknown-unknown -mode=latency -snippets-file=%s | FileCheck %s + +# CHECK: measurements: + +# LLVM-EXEGESIS-DEFREG R11 ff +# LLVM-EXEGESIS-LOOP-REGISTER R12 + +addq $0xff, %r11 diff --git a/llvm/test/tools/llvm-objdump/X86/elf-pgoanalysismap.yaml b/llvm/test/tools/llvm-objdump/X86/elf-pgoanalysismap.yaml index 732fab3e2a37..4d1e5408d86d 100644 --- a/llvm/test/tools/llvm-objdump/X86/elf-pgoanalysismap.yaml +++ b/llvm/test/tools/llvm-objdump/X86/elf-pgoanalysismap.yaml @@ -47,7 +47,9 @@ Symbols: # RUN: yaml2obj %s --docnum=2 -o %t2 # RUN: llvm-objdump %t2 -d --symbolize-operands --no-show-raw-insn --no-leading-addr | \ -# RUN: FileCheck %s --check-prefix=ENTRYCOUNT-BLOCKFREQ +# RUN: FileCheck --match-full-lines --strict-whitespace %s --check-prefix=ENTRYCOUNT-BLOCKFREQ +# RUN: llvm-objdump %t2 -d --symbolize-operands --pretty-pgo-analysis-map --no-show-raw-insn --no-leading-addr | \ +# RUN: FileCheck --match-full-lines --strict-whitespace %s --check-prefix=ENTRYCOUNT-BLOCKFREQ-PRETTY --- !ELF FileHeader: @@ -98,18 +100,28 @@ Symbols: Section: .text.foo Value: 0x0 -# ENTRYCOUNT-BLOCKFREQ: <foo>: -# ENTRYCOUNT-BLOCKFREQ: <BB3> (Entry count: 1000, Frequency: 1000): -# ENTRYCOUNT-BLOCKFREQ: <BB1> (Frequency: 133): -# ENTRYCOUNT-BLOCKFREQ: <BB2> (Frequency: 18): -# ENTRYCOUNT-BLOCKFREQ: <BB5> (Frequency: 1000): +# ENTRYCOUNT-BLOCKFREQ:<foo>: +# ENTRYCOUNT-BLOCKFREQ:<BB3> (Entry count: 1000, Frequency: 1000): +# ENTRYCOUNT-BLOCKFREQ:<BB1> (Frequency: 133): +# ENTRYCOUNT-BLOCKFREQ:<BB2> (Frequency: 18): +# ENTRYCOUNT-BLOCKFREQ:<BB5> (Frequency: 1000): + +# ENTRYCOUNT-BLOCKFREQ-PRETTY:<foo>: +# ENTRYCOUNT-BLOCKFREQ-PRETTY:<BB3> (Entry count: 1000, Frequency: 1.0): +# ENTRYCOUNT-BLOCKFREQ-PRETTY:<BB1> (Frequency: 0.133): +# ENTRYCOUNT-BLOCKFREQ-PRETTY:<BB2> (Frequency: 0.018): +# ENTRYCOUNT-BLOCKFREQ-PRETTY:<BB5> (Frequency: 1.0): ## Check the case where we have entry points, block frequency, and branch ## proabability information. # RUN: yaml2obj %s --docnum=3 -o %t3 # RUN: llvm-objdump %t3 -d --symbolize-operands --no-show-raw-insn --no-leading-addr | \ -# RUN: FileCheck %s --check-prefix=ENTRY-FREQ-PROB +# RUN: FileCheck --match-full-lines --strict-whitespace %s --check-prefix=ENTRY-FREQ-PROB +# RUN: llvm-objdump %t3 -d --symbolize-operands --pretty-pgo-analysis-map --no-show-raw-insn --no-leading-addr | \ +# RUN: FileCheck --match-full-lines --strict-whitespace %s --check-prefix=ENTRY-FREQ-PROB-PRETTY +# RUN: llvm-objdump %t3 -d --pretty-pgo-analysis-map --no-show-raw-insn --no-leading-addr 2>&1 | \ +# RUN: FileCheck %s --check-prefix=MISSING-SYMBOLIZE-OPERANDS --- !ELF FileHeader: @@ -154,21 +166,21 @@ Sections: - BBFreq: 1000 Successors: - ID: 1 - BrProb: 0x22222222 + BrProb: 0x10000000 - ID: 2 - BrProb: 0x33333333 + BrProb: 0x15000000 - ID: 3 - BrProb: 0xaaaaaaaa + BrProb: 0x50000000 - BBFreq: 133 Successors: - ID: 2 - BrProb: 0x11111111 + BrProb: 0x10000000 - ID: 3 - BrProb: 0xeeeeeeee + BrProb: 0x70000000 - BBFreq: 18 Successors: - ID: 3 - BrProb: 0xffffffff + BrProb: 0x80000000 - BBFreq: 1000 Successors: [] Symbols: @@ -176,8 +188,16 @@ Symbols: Section: .text.foo Value: 0x0 -# ENTRY-FREQ-PROB: <foo>: -# ENTRY-FREQ-PROB: <BB3> (Entry count: 1000, Frequency: 1000, Successors: BB1:22222222, BB2:33333333, BB3:aaaaaaaa): -# ENTRY-FREQ-PROB: <BB1> (Frequency: 133, Successors: BB2:11111111, BB3:eeeeeeee): -# ENTRY-FREQ-PROB: <BB2> (Frequency: 18, Successors: BB3:ffffffff): -# ENTRY-FREQ-PROB: <BB5> (Frequency: 1000): +# ENTRY-FREQ-PROB:<foo>: +# ENTRY-FREQ-PROB:<BB3> (Entry count: 1000, Frequency: 1000, Successors: BB1:10000000, BB2:15000000, BB3:50000000): +# ENTRY-FREQ-PROB:<BB1> (Frequency: 133, Successors: BB2:10000000, BB3:70000000): +# ENTRY-FREQ-PROB:<BB2> (Frequency: 18, Successors: BB3:80000000): +# ENTRY-FREQ-PROB:<BB5> (Frequency: 1000): + +# ENTRY-FREQ-PROB-PRETTY:<foo>: +# ENTRY-FREQ-PROB-PRETTY:<BB3> (Entry count: 1000, Frequency: 1.0, Successors: BB1:[0x10000000 / 0x80000000 = 12.50%], BB2:[0x15000000 / 0x80000000 = 16.41%], BB3:[0x50000000 / 0x80000000 = 62.50%]): +# ENTRY-FREQ-PROB-PRETTY:<BB1> (Frequency: 0.133, Successors: BB2:[0x10000000 / 0x80000000 = 12.50%], BB3:[0x70000000 / 0x80000000 = 87.50%]): +# ENTRY-FREQ-PROB-PRETTY:<BB2> (Frequency: 0.018, Successors: BB3:[0x80000000 / 0x80000000 = 100.00%]): +# ENTRY-FREQ-PROB-PRETTY:<BB5> (Frequency: 1.0): + +# MISSING-SYMBOLIZE-OPERANDS: warning: --symbolize-operands must be enabled for --pretty-pgo-analysis-map to have an effect diff --git a/llvm/test/tools/llvm-readobj/ELF/bb-addr-map-pgo-analysis-map.test b/llvm/test/tools/llvm-readobj/ELF/bb-addr-map-pgo-analysis-map.test index e5a9400c670c..5faafd4d83b2 100644 --- a/llvm/test/tools/llvm-readobj/ELF/bb-addr-map-pgo-analysis-map.test +++ b/llvm/test/tools/llvm-readobj/ELF/bb-addr-map-pgo-analysis-map.test @@ -3,17 +3,19 @@ ## Check 64-bit: # RUN: yaml2obj %s -DBITS=64 -DADDR=0x999999999 -o %t1.x64.o -# RUN: llvm-readobj %t1.x64.o --bb-addr-map 2>&1 | FileCheck %s -DADDR=0x999999999 -DFILE=%t1.x64.o --check-prefix=CHECK +# RUN: llvm-readobj %t1.x64.o --bb-addr-map 2>&1 | FileCheck --match-full-lines %s -DADDR=0x999999999 -DFILE=%t1.x64.o --check-prefixes=CHECK,RAW +# RUN: llvm-readobj %t1.x64.o --bb-addr-map --pretty-pgo-analysis-map 2>&1 | FileCheck --match-full-lines %s -DADDR=0x999999999 -DFILE=%t1.x64.o --check-prefixes=CHECK,PRETTY # RUN: llvm-readelf %t1.x64.o --bb-addr-map | FileCheck %s --check-prefix=GNU +# RUN: llvm-readobj %t1.x64.o --pretty-pgo-analysis-map 2>&1 | FileCheck %s --check-prefix=PRETTY-NO-BAM ## Check 32-bit: # RUN: yaml2obj %s -DBITS=32 -o %t1.x32.o -# RUN: llvm-readobj %t1.x32.o --bb-addr-map 2>&1 | FileCheck -DADDR=0x11111 %s -DFILE=%t1.x32.o --check-prefix=CHECK +# RUN: llvm-readobj %t1.x32.o --bb-addr-map 2>&1 | FileCheck --match-full-lines -DADDR=0x11111 %s -DFILE=%t1.x32.o --check-prefixes=CHECK,RAW # RUN: llvm-readelf %t1.x32.o --bb-addr-map | FileCheck %s --check-prefix=GNU ## Check that a malformed section can be handled. # RUN: yaml2obj %s -DBITS=32 -DSIZE=24 -o %t2.o -# RUN: llvm-readobj %t2.o --bb-addr-map 2>&1 | FileCheck %s -DOFFSET=0x00000018 -DFILE=%t2.o --check-prefix=TRUNCATED +# RUN: llvm-readobj %t2.o --bb-addr-map 2>&1 | FileCheck --match-full-lines %s -DOFFSET=0x00000018 -DFILE=%t2.o --check-prefix=TRUNCATED ## Check that missing features can be handled. # RUN: yaml2obj %s -DBITS=32 -DFEATURE=0x2 -o %t3.o @@ -22,7 +24,7 @@ # CHECK: BBAddrMap [ # CHECK-NEXT: Function { # CHECK-NEXT: At: [[ADDR]] -# CHECK-NEXT: warning: '[[FILE]]': could not identify function symbol for address ([[ADDR]]) in SHT_LLVM_BB_ADDR_MAP section with index 3 +# CHECK-NEXT: {{.*}}: warning: '[[FILE]]': could not identify function symbol for address ([[ADDR]]) in SHT_LLVM_BB_ADDR_MAP section with index 3 # CHECK-NEXT: Name: <?> # CHECK-NEXT: BB Ranges [ # CHECK-NEXT: { @@ -55,16 +57,19 @@ # CHECK-NEXT: FuncEntryCount: 100 # CHECK-NEXT: PGO BB entries [ # CHECK-NEXT: { -# CHECK-NEXT: Frequency: 100 +# RAW-NEXT: Frequency: 100 +# PRETTY-NEXT: Frequency: 1.0 # CHECK-NEXT: Successors [ # CHECK-NEXT: { # CHECK-NEXT: ID: 2 -# CHECK-NEXT: Probability: 0xFFFFFFFF +# RAW-NEXT: Probability: 0x80000000 +# PRETTY-NEXT: Probability: 0x80000000 / 0x80000000 = 100.00% # CHECK-NEXT: } # CHECK-NEXT: ] # CHECK-NEXT: } # CHECK-NEXT: { -# CHECK-NEXT: Frequency: 100 +# RAW-NEXT: Frequency: 100 +# PRETTY-NEXT: Frequency: 1.0 # CHECK-NEXT: Successors [ # CHECK-NEXT: ] # CHECK-NEXT: } @@ -95,7 +100,8 @@ # CHECK-NEXT: FuncEntryCount: 8888 # CHECK-NEXT: PGO BB entries [ # CHECK-NEXT: { -# CHECK-NEXT: Frequency: 9000 +# RAW-NEXT: Frequency: 9000 +# PRETTY-NEXT: Frequency: 1.0 # CHECK-NEXT: } # CHECK-NEXT: ] # CHECK-NEXT: } @@ -104,8 +110,10 @@ # GNU: GNUStyle::printBBAddrMaps not implemented +# PRETTY-NO-BAM: warning: --bb-addr-map must be enabled for --pretty-pgo-analysis-map to have an effect + # TRUNCATED: BBAddrMap [ -# TRUNCATED-NEXT: warning: '[[FILE]]': unable to dump SHT_LLVM_BB_ADDR_MAP section with index 3: unable to decode LEB128 at offset [[OFFSET]]: malformed uleb128, extends past end +# TRUNCATED-NEXT: {{.*}}: warning: '[[FILE]]': unable to dump SHT_LLVM_BB_ADDR_MAP section with index 3: unable to decode LEB128 at offset [[OFFSET]]: malformed uleb128, extends past end # TRUNCATED-NEXT: ] ## Check that the other valid section is properly dumped. # TRUNCATED-NEXT: BBAddrMap [ @@ -192,7 +200,7 @@ Sections: - BBFreq: 100 Successors: - ID: 2 - BrProb: 0xFFFFFFFF + BrProb: 0x80000000 - BBFreq: 100 Successors: [] - FuncEntryCount: 8888 diff --git a/llvm/tools/llvm-exegesis/lib/BenchmarkResult.h b/llvm/tools/llvm-exegesis/lib/BenchmarkResult.h index 0aecaaeea4b2..4ae6bc2a54cd 100644 --- a/llvm/tools/llvm-exegesis/lib/BenchmarkResult.h +++ b/llvm/tools/llvm-exegesis/lib/BenchmarkResult.h @@ -74,6 +74,8 @@ struct BenchmarkKey { // The address that the snippet should be loaded in at if the execution mode // being used supports it. intptr_t SnippetAddress = 0; + // The register that should be used to hold the loop counter. + unsigned LoopRegister; }; struct BenchmarkMeasure { diff --git a/llvm/tools/llvm-exegesis/lib/SnippetFile.cpp b/llvm/tools/llvm-exegesis/lib/SnippetFile.cpp index 7258fcb4279c..431d99c72b80 100644 --- a/llvm/tools/llvm-exegesis/lib/SnippetFile.cpp +++ b/llvm/tools/llvm-exegesis/lib/SnippetFile.cpp @@ -9,6 +9,7 @@ #include "SnippetFile.h" #include "BenchmarkRunner.h" #include "Error.h" +#include "Target.h" #include "llvm/MC/MCContext.h" #include "llvm/MC/MCInstPrinter.h" #include "llvm/MC/MCObjectFileInfo.h" @@ -175,6 +176,20 @@ public: return; } + if (CommentText.consume_front("LOOP-REGISTER")) { + // LLVM-EXEGESIS-LOOP-REGISTER <loop register> + unsigned LoopRegister; + + if (!(LoopRegister = findRegisterByName(CommentText.trim()))) { + errs() << "unknown register '" << CommentText + << "' in 'LLVM-EXEGESIS-LOOP-REGISTER " << CommentText << "'\n"; + ++InvalidComments; + return; + } + + Result->Key.LoopRegister = LoopRegister; + return; + } } unsigned numInvalidComments() const { return InvalidComments; } @@ -221,6 +236,11 @@ Expected<std::vector<BenchmarkCode>> readSnippets(const LLVMState &State, BenchmarkCode Result; + // Ensure that there is a default loop register value specified. + Result.Key.LoopRegister = + State.getExegesisTarget().getDefaultLoopCounterRegister( + State.getTargetMachine().getTargetTriple()); + const TargetMachine &TM = State.getTargetMachine(); MCContext Context(TM.getTargetTriple(), TM.getMCAsmInfo(), TM.getMCRegisterInfo(), TM.getMCSubtargetInfo()); diff --git a/llvm/tools/llvm-exegesis/lib/SnippetRepetitor.cpp b/llvm/tools/llvm-exegesis/lib/SnippetRepetitor.cpp index 561687a62319..0bab30d15820 100644 --- a/llvm/tools/llvm-exegesis/lib/SnippetRepetitor.cpp +++ b/llvm/tools/llvm-exegesis/lib/SnippetRepetitor.cpp @@ -48,10 +48,8 @@ public: class LoopSnippetRepetitor : public SnippetRepetitor { public: - explicit LoopSnippetRepetitor(const LLVMState &State) - : SnippetRepetitor(State), - LoopCounter(State.getExegesisTarget().getLoopCounterRegister( - State.getTargetMachine().getTargetTriple())) {} + explicit LoopSnippetRepetitor(const LLVMState &State, unsigned LoopRegister) + : SnippetRepetitor(State), LoopCounter(LoopRegister) {} // Loop over the snippet ceil(MinInstructions / Instructions.Size()) times. FillFunction Repeat(ArrayRef<MCInst> Instructions, unsigned MinInstructions, @@ -113,8 +111,8 @@ public: (void)_; Loop.addInstructions(Instructions); } - ET.decrementLoopCounterAndJump(*Loop.MBB, *Loop.MBB, - State.getInstrInfo()); + ET.decrementLoopCounterAndJump(*Loop.MBB, *Loop.MBB, State.getInstrInfo(), + LoopCounter); // Set up the exit basic block. Loop.MBB->addSuccessor(Exit.MBB, BranchProbability::getZero()); @@ -138,14 +136,14 @@ SnippetRepetitor::~SnippetRepetitor() {} std::unique_ptr<const SnippetRepetitor> SnippetRepetitor::Create(Benchmark::RepetitionModeE Mode, - const LLVMState &State) { + const LLVMState &State, unsigned LoopRegister) { switch (Mode) { case Benchmark::Duplicate: case Benchmark::MiddleHalfDuplicate: return std::make_unique<DuplicateSnippetRepetitor>(State); case Benchmark::Loop: case Benchmark::MiddleHalfLoop: - return std::make_unique<LoopSnippetRepetitor>(State); + return std::make_unique<LoopSnippetRepetitor>(State, LoopRegister); case Benchmark::AggregateMin: break; } diff --git a/llvm/tools/llvm-exegesis/lib/SnippetRepetitor.h b/llvm/tools/llvm-exegesis/lib/SnippetRepetitor.h index 2b3c416c9029..c62e80f161f1 100644 --- a/llvm/tools/llvm-exegesis/lib/SnippetRepetitor.h +++ b/llvm/tools/llvm-exegesis/lib/SnippetRepetitor.h @@ -29,7 +29,8 @@ namespace exegesis { class SnippetRepetitor { public: static std::unique_ptr<const SnippetRepetitor> - Create(Benchmark::RepetitionModeE Mode, const LLVMState &State); + Create(Benchmark::RepetitionModeE Mode, const LLVMState &State, + unsigned LoopRegister); virtual ~SnippetRepetitor(); diff --git a/llvm/tools/llvm-exegesis/lib/Target.h b/llvm/tools/llvm-exegesis/lib/Target.h index 7bbd946b0333..522c75d15703 100644 --- a/llvm/tools/llvm-exegesis/lib/Target.h +++ b/llvm/tools/llvm-exegesis/lib/Target.h @@ -202,12 +202,15 @@ public: } // Returns a counter usable as a loop counter. - virtual unsigned getLoopCounterRegister(const Triple &) const { return 0; } + virtual unsigned getDefaultLoopCounterRegister(const Triple &) const { + return 0; + } // Adds the code to decrement the loop counter and virtual void decrementLoopCounterAndJump(MachineBasicBlock &MBB, MachineBasicBlock &TargetMBB, - const MCInstrInfo &MII) const { + const MCInstrInfo &MII, + unsigned LoopRegister) const { llvm_unreachable("decrementLoopCounterAndBranch() requires " "getLoopCounterRegister() > 0"); } diff --git a/llvm/tools/llvm-exegesis/lib/X86/Target.cpp b/llvm/tools/llvm-exegesis/lib/X86/Target.cpp index 6fc951a6e35d..a41a995f5560 100644 --- a/llvm/tools/llvm-exegesis/lib/X86/Target.cpp +++ b/llvm/tools/llvm-exegesis/lib/X86/Target.cpp @@ -720,7 +720,7 @@ private: unsigned getScratchMemoryRegister(const Triple &TT) const override; - unsigned getLoopCounterRegister(const Triple &) const override; + unsigned getDefaultLoopCounterRegister(const Triple &) const override; unsigned getMaxMemoryAccessSize() const override { return 64; } @@ -733,7 +733,8 @@ private: void decrementLoopCounterAndJump(MachineBasicBlock &MBB, MachineBasicBlock &TargetMBB, - const MCInstrInfo &MII) const override; + const MCInstrInfo &MII, + unsigned LoopRegister) const override; std::vector<MCInst> setRegTo(const MCSubtargetInfo &STI, unsigned Reg, const APInt &Value) const override; @@ -852,7 +853,7 @@ const unsigned ExegesisX86Target::kUnavailableRegistersSSE[12] = { // We're using one of R8-R15 because these registers are never hardcoded in // instructions (e.g. MOVS writes to EDI, ESI, EDX), so they have less // conflicts. -constexpr const unsigned kLoopCounterReg = X86::R8; +constexpr const unsigned kDefaultLoopCounterReg = X86::R8; } // namespace @@ -870,11 +871,12 @@ unsigned ExegesisX86Target::getScratchMemoryRegister(const Triple &TT) const { return TT.isOSWindows() ? X86::RCX : X86::RDI; } -unsigned ExegesisX86Target::getLoopCounterRegister(const Triple &TT) const { +unsigned +ExegesisX86Target::getDefaultLoopCounterRegister(const Triple &TT) const { if (!TT.isArch64Bit()) { return 0; } - return kLoopCounterReg; + return kDefaultLoopCounterReg; } Error ExegesisX86Target::randomizeTargetMCOperand( @@ -912,10 +914,10 @@ void ExegesisX86Target::fillMemoryOperands(InstructionTemplate &IT, void ExegesisX86Target::decrementLoopCounterAndJump( MachineBasicBlock &MBB, MachineBasicBlock &TargetMBB, - const MCInstrInfo &MII) const { + const MCInstrInfo &MII, unsigned LoopRegister) const { BuildMI(&MBB, DebugLoc(), MII.get(X86::ADD64ri8)) - .addDef(kLoopCounterReg) - .addUse(kLoopCounterReg) + .addDef(LoopRegister) + .addUse(LoopRegister) .addImm(-1); BuildMI(&MBB, DebugLoc(), MII.get(X86::JCC_1)) .addMBB(&TargetMBB) diff --git a/llvm/tools/llvm-exegesis/llvm-exegesis.cpp b/llvm/tools/llvm-exegesis/llvm-exegesis.cpp index 782d44422791..1ae2565e894c 100644 --- a/llvm/tools/llvm-exegesis/llvm-exegesis.cpp +++ b/llvm/tools/llvm-exegesis/llvm-exegesis.cpp @@ -497,22 +497,42 @@ void benchmarkMain() { } const auto Opcodes = getOpcodesOrDie(State); + std::vector<BenchmarkCode> Configurations; + + unsigned LoopRegister = + State.getExegesisTarget().getDefaultLoopCounterRegister( + State.getTargetMachine().getTargetTriple()); + + if (Opcodes.empty()) { + Configurations = ExitOnErr(readSnippets(State, SnippetsFile)); + for (const auto &Configuration : Configurations) { + if (ExecutionMode != BenchmarkRunner::ExecutionModeE::SubProcess && + (Configuration.Key.MemoryMappings.size() != 0 || + Configuration.Key.MemoryValues.size() != 0 || + Configuration.Key.SnippetAddress != 0)) + ExitWithError("Memory and snippet address annotations are only " + "supported in subprocess " + "execution mode"); + } + LoopRegister = Configurations[0].Key.LoopRegister; + } SmallVector<std::unique_ptr<const SnippetRepetitor>, 2> Repetitors; if (RepetitionMode != Benchmark::RepetitionModeE::AggregateMin) - Repetitors.emplace_back(SnippetRepetitor::Create(RepetitionMode, State)); + Repetitors.emplace_back( + SnippetRepetitor::Create(RepetitionMode, State, LoopRegister)); else { for (Benchmark::RepetitionModeE RepMode : {Benchmark::RepetitionModeE::Duplicate, Benchmark::RepetitionModeE::Loop}) - Repetitors.emplace_back(SnippetRepetitor::Create(RepMode, State)); + Repetitors.emplace_back( + SnippetRepetitor::Create(RepMode, State, LoopRegister)); } BitVector AllReservedRegs; for (const std::unique_ptr<const SnippetRepetitor> &Repetitor : Repetitors) AllReservedRegs |= Repetitor->getReservedRegs(); - std::vector<BenchmarkCode> Configurations; if (!Opcodes.empty()) { for (const unsigned Opcode : Opcodes) { // Ignore instructions without a sched class if @@ -534,17 +554,6 @@ void benchmarkMain() { std::move(ConfigsForInstr->begin(), ConfigsForInstr->end(), std::back_inserter(Configurations)); } - } else { - Configurations = ExitOnErr(readSnippets(State, SnippetsFile)); - for (const auto &Configuration : Configurations) { - if (ExecutionMode != BenchmarkRunner::ExecutionModeE::SubProcess && - (Configuration.Key.MemoryMappings.size() != 0 || - Configuration.Key.MemoryValues.size() != 0 || - Configuration.Key.SnippetAddress != 0)) - ExitWithError("Memory and snippet address annotations are only " - "supported in subprocess " - "execution mode"); - } } if (MinInstructions == 0) { diff --git a/llvm/tools/llvm-objdump/ObjdumpOpts.td b/llvm/tools/llvm-objdump/ObjdumpOpts.td index c1dec5ced89d..c3764c6e9753 100644 --- a/llvm/tools/llvm-objdump/ObjdumpOpts.td +++ b/llvm/tools/llvm-objdump/ObjdumpOpts.td @@ -210,6 +210,10 @@ def : Flag<["-"], "t">, Alias<syms>, HelpText<"Alias for --syms">; def symbolize_operands : Flag<["--"], "symbolize-operands">, HelpText<"Symbolize instruction operands when disassembling">; +def pretty_pgo_analysis_map : Flag<["--"], "pretty-pgo-analysis-map">, + HelpText<"Display PGO analysis values with " + "formatting rather than raw numbers">; + def dynamic_syms : Flag<["--"], "dynamic-syms">, HelpText<"Display the contents of the dynamic symbol table">; def : Flag<["-"], "T">, Alias<dynamic_syms>, diff --git a/llvm/tools/llvm-objdump/llvm-objdump.cpp b/llvm/tools/llvm-objdump/llvm-objdump.cpp index 948a5d74e1ab..78cf67b1e630 100644 --- a/llvm/tools/llvm-objdump/llvm-objdump.cpp +++ b/llvm/tools/llvm-objdump/llvm-objdump.cpp @@ -188,8 +188,10 @@ public: const BBAddrMap &getAddrMap() const { return AddrMap; } // Returns the PGO string associated with the entry of index `PGOBBEntryIndex` - // in `PGOMap`. - std::string constructPGOLabelString(size_t PGOBBEntryIndex) const { + // in `PGOMap`. If PrettyPGOAnalysis is true, prints BFI as relative frequency + // and BPI as percentage. Otherwise raw values are displayed. + std::string constructPGOLabelString(size_t PGOBBEntryIndex, + bool PrettyPGOAnalysis) const { if (!PGOMap.FeatEnable.hasPGOAnalysis()) return ""; std::string PGOString; @@ -211,7 +213,12 @@ public: PGOMap.BBEntries[PGOBBEntryIndex]; if (PGOMap.FeatEnable.BBFreq) { - PGOSS << "Frequency: " << Twine(PGOBBEntry.BlockFreq.getFrequency()); + PGOSS << "Frequency: "; + if (PrettyPGOAnalysis) + printRelativeBlockFreq(PGOSS, PGOMap.BBEntries.front().BlockFreq, + PGOBBEntry.BlockFreq); + else + PGOSS << Twine(PGOBBEntry.BlockFreq.getFrequency()); if (PGOMap.FeatEnable.BrProb && PGOBBEntry.Successors.size() > 0) { PGOSS << ", "; } @@ -220,9 +227,12 @@ public: PGOSS << "Successors: "; interleaveComma( PGOBBEntry.Successors, PGOSS, - [&PGOSS](const PGOAnalysisMap::PGOBBEntry::SuccessorEntry &SE) { + [&](const PGOAnalysisMap::PGOBBEntry::SuccessorEntry &SE) { PGOSS << "BB" << SE.ID << ":"; - PGOSS.write_hex(SE.Prob.getNumerator()); + if (PrettyPGOAnalysis) + PGOSS << "[" << SE.Prob << "]"; + else + PGOSS.write_hex(SE.Prob.getNumerator()); }); } } @@ -331,6 +341,7 @@ static bool HasStopAddressFlag; bool objdump::SymbolTable; static bool SymbolizeOperands; +static bool PrettyPGOAnalysisMap; static bool DynamicSymbolTable; std::string objdump::TripleName; bool objdump::UnwindInfo; @@ -1410,8 +1421,8 @@ static void collectBBAddrMapLabels( std::string LabelString = ("BB" + Twine(BBEntry.ID)).str(); Labels[BBAddress].push_back( - {LabelString, - FunctionMap->constructPGOLabelString(NumBBEntriesBeforeRange + I)}); + {LabelString, FunctionMap->constructPGOLabelString( + NumBBEntriesBeforeRange + I, PrettyPGOAnalysisMap)}); } } @@ -3473,6 +3484,10 @@ static void parseObjdumpOptions(const llvm::opt::InputArgList &InputArgs) { HasStopAddressFlag = InputArgs.hasArg(OBJDUMP_stop_address_EQ); SymbolTable = InputArgs.hasArg(OBJDUMP_syms); SymbolizeOperands = InputArgs.hasArg(OBJDUMP_symbolize_operands); + PrettyPGOAnalysisMap = InputArgs.hasArg(OBJDUMP_pretty_pgo_analysis_map); + if (PrettyPGOAnalysisMap && !SymbolizeOperands) + reportCmdLineWarning("--symbolize-operands must be enabled for " + "--pretty-pgo-analysis-map to have an effect"); DynamicSymbolTable = InputArgs.hasArg(OBJDUMP_dynamic_syms); TripleName = InputArgs.getLastArgValue(OBJDUMP_triple_EQ).str(); UnwindInfo = InputArgs.hasArg(OBJDUMP_unwind_info); diff --git a/llvm/tools/llvm-readobj/ELFDumper.cpp b/llvm/tools/llvm-readobj/ELFDumper.cpp index 4be678df4412..e78732353cc8 100644 --- a/llvm/tools/llvm-readobj/ELFDumper.cpp +++ b/llvm/tools/llvm-readobj/ELFDumper.cpp @@ -593,7 +593,7 @@ public: void printVersionDefinitionSection(const Elf_Shdr *Sec) override; void printVersionDependencySection(const Elf_Shdr *Sec) override; void printCGProfile() override; - void printBBAddrMaps() override; + void printBBAddrMaps(bool PrettyPGOAnalysis) override; void printAddrsig() override; void printNotes() override; void printELFLinkerOptions() override; @@ -704,7 +704,7 @@ public: void printVersionDefinitionSection(const Elf_Shdr *Sec) override; void printVersionDependencySection(const Elf_Shdr *Sec) override; void printCGProfile() override; - void printBBAddrMaps() override; + void printBBAddrMaps(bool PrettyPGOAnalysis) override; void printAddrsig() override; void printNotes() override; void printELFLinkerOptions() override; @@ -5036,7 +5036,8 @@ template <class ELFT> void GNUELFDumper<ELFT>::printCGProfile() { OS << "GNUStyle::printCGProfile not implemented\n"; } -template <class ELFT> void GNUELFDumper<ELFT>::printBBAddrMaps() { +template <class ELFT> +void GNUELFDumper<ELFT>::printBBAddrMaps(bool /*PrettyPGOAnalysis*/) { OS << "GNUStyle::printBBAddrMaps not implemented\n"; } @@ -7526,7 +7527,8 @@ template <class ELFT> void LLVMELFDumper<ELFT>::printCGProfile() { } } -template <class ELFT> void LLVMELFDumper<ELFT>::printBBAddrMaps() { +template <class ELFT> +void LLVMELFDumper<ELFT>::printBBAddrMaps(bool PrettyPGOAnalysis) { bool IsRelocatable = this->Obj.getHeader().e_type == ELF::ET_REL; using Elf_Shdr = typename ELFT::Shdr; auto IsMatch = [](const Elf_Shdr &Sec) -> bool { @@ -7605,21 +7607,28 @@ template <class ELFT> void LLVMELFDumper<ELFT>::printBBAddrMaps() { for (const PGOAnalysisMap::PGOBBEntry &PBBE : PAM.BBEntries) { DictScope L(W); - /// FIXME: currently we just emit the raw frequency, it may be - /// better to provide an option to scale it by the first entry - /// frequence using BlockFrequency::Scaled64 number - if (PAM.FeatEnable.BBFreq) - W.printNumber("Frequency", PBBE.BlockFreq.getFrequency()); + if (PAM.FeatEnable.BBFreq) { + if (PrettyPGOAnalysis) { + std::string BlockFreqStr; + raw_string_ostream SS(BlockFreqStr); + printRelativeBlockFreq(SS, PAM.BBEntries.front().BlockFreq, + PBBE.BlockFreq); + W.printString("Frequency", BlockFreqStr); + } else { + W.printNumber("Frequency", PBBE.BlockFreq.getFrequency()); + } + } if (PAM.FeatEnable.BrProb) { ListScope L(W, "Successors"); for (const auto &Succ : PBBE.Successors) { DictScope L(W); W.printNumber("ID", Succ.ID); - /// FIXME: currently we just emit the raw numerator of the - /// probably, it may be better to provide an option to emit it - /// as a percentage or other prettied representation - W.printHex("Probability", Succ.Prob.getNumerator()); + if (PrettyPGOAnalysis) { + W.printObject("Probability", Succ.Prob); + } else { + W.printHex("Probability", Succ.Prob.getNumerator()); + } } } } diff --git a/llvm/tools/llvm-readobj/ObjDumper.h b/llvm/tools/llvm-readobj/ObjDumper.h index 3958dd3a3333..cd744e3bbfb7 100644 --- a/llvm/tools/llvm-readobj/ObjDumper.h +++ b/llvm/tools/llvm-readobj/ObjDumper.h @@ -129,7 +129,9 @@ public: virtual void printGroupSections() {} virtual void printHashHistograms() {} virtual void printCGProfile() {} - virtual void printBBAddrMaps() {} + // If PrettyPGOAnalysis is true, prints BFI as relative frequency and BPI as + // percentage. Otherwise raw values are displayed. + virtual void printBBAddrMaps(bool PrettyPGOAnalysis) {} virtual void printAddrsig() {} virtual void printNotes() {} virtual void printELFLinkerOptions() {} diff --git a/llvm/tools/llvm-readobj/Opts.td b/llvm/tools/llvm-readobj/Opts.td index 018facc278e8..1e9cde6b2e87 100644 --- a/llvm/tools/llvm-readobj/Opts.td +++ b/llvm/tools/llvm-readobj/Opts.td @@ -19,6 +19,7 @@ def all : FF<"all", "Equivalent to setting: --file-header, --program-headers, -- "--section-groups and --histogram">; def arch_specific : FF<"arch-specific", "Display architecture-specific information">; def bb_addr_map : FF<"bb-addr-map", "Display the BB address map section">; +def pretty_pgo_analysis_map : FF<"pretty-pgo-analysis-map", "Display PGO analysis values with formatting rather than raw numbers">; def cg_profile : FF<"cg-profile", "Display call graph profile section">; def decompress : FF<"decompress", "Dump decompressed section content when used with -x or -p">; defm demangle : BB<"demangle", "Demangle symbol names", "Do not demangle symbol names (default)">; diff --git a/llvm/tools/llvm-readobj/llvm-readobj.cpp b/llvm/tools/llvm-readobj/llvm-readobj.cpp index 979433d69011..a0b576566016 100644 --- a/llvm/tools/llvm-readobj/llvm-readobj.cpp +++ b/llvm/tools/llvm-readobj/llvm-readobj.cpp @@ -95,6 +95,7 @@ static bool Addrsig; static bool All; static bool ArchSpecificInfo; static bool BBAddrMap; +static bool PrettyPGOAnalysisMap; bool ExpandRelocs; static bool CGProfile; static bool Decompress; @@ -212,6 +213,11 @@ static void parseOptions(const opt::InputArgList &Args) { opts::All = Args.hasArg(OPT_all); opts::ArchSpecificInfo = Args.hasArg(OPT_arch_specific); opts::BBAddrMap = Args.hasArg(OPT_bb_addr_map); + opts::PrettyPGOAnalysisMap = Args.hasArg(OPT_pretty_pgo_analysis_map); + if (opts::PrettyPGOAnalysisMap && !opts::BBAddrMap) + WithColor::warning(errs(), ToolName) + << "--bb-addr-map must be enabled for --pretty-pgo-analysis-map to " + "have an effect\n"; opts::CGProfile = Args.hasArg(OPT_cg_profile); opts::Decompress = Args.hasArg(OPT_decompress); opts::Demangle = Args.hasFlag(OPT_demangle, OPT_no_demangle, false); @@ -466,7 +472,7 @@ static void dumpObject(ObjectFile &Obj, ScopedPrinter &Writer, if (opts::CGProfile) Dumper->printCGProfile(); if (opts::BBAddrMap) - Dumper->printBBAddrMaps(); + Dumper->printBBAddrMaps(opts::PrettyPGOAnalysisMap); if (opts::Addrsig) Dumper->printAddrsig(); if (opts::Notes) diff --git a/llvm/unittests/tools/llvm-exegesis/X86/SnippetFileTest.cpp b/llvm/unittests/tools/llvm-exegesis/X86/SnippetFileTest.cpp index 505a030675f6..f1fa89117117 100644 --- a/llvm/unittests/tools/llvm-exegesis/X86/SnippetFileTest.cpp +++ b/llvm/unittests/tools/llvm-exegesis/X86/SnippetFileTest.cpp @@ -219,6 +219,25 @@ TEST_F(X86SnippetFileTest, SnippetAddress) { EXPECT_EQ(Snippet.Key.SnippetAddress, 0x10000); } +TEST_F(X86SnippetFileTest, LoopRegister) { + auto Snippets = TestCommon(R"( + # LLVM-EXEGESIS-LOOP-REGISTER R11 + )"); + ASSERT_TRUE(static_cast<bool>(Snippets)); + EXPECT_THAT(*Snippets, SizeIs(1)); + const auto &Snippet = (*Snippets)[0]; + EXPECT_EQ(Snippet.Key.LoopRegister, X86::R11); +} + +TEST_F(X86SnippetFileTest, LoopRegisterInvalidRegister) { + auto Error = TestCommon(R"( + # LLVM-EXEGESIS-LOOP-REGISTER INVALID + )") + .takeError(); + EXPECT_TRUE(static_cast<bool>(Error)); + consumeError(std::move(Error)); +} + } // namespace } // namespace exegesis } // namespace llvm diff --git a/llvm/unittests/tools/llvm-exegesis/X86/SnippetRepetitorTest.cpp b/llvm/unittests/tools/llvm-exegesis/X86/SnippetRepetitorTest.cpp index 25e8836087c1..b55ca5057ae0 100644 --- a/llvm/unittests/tools/llvm-exegesis/X86/SnippetRepetitorTest.cpp +++ b/llvm/unittests/tools/llvm-exegesis/X86/SnippetRepetitorTest.cpp @@ -40,7 +40,10 @@ protected: void TestCommon(Benchmark::RepetitionModeE RepetitionMode, unsigned SnippetInstructions = 1) { - const auto Repetitor = SnippetRepetitor::Create(RepetitionMode, State); + const auto Repetitor = SnippetRepetitor::Create( + RepetitionMode, State, + State.getExegesisTarget().getDefaultLoopCounterRegister( + State.getTargetMachine().getTargetTriple())); const std::vector<MCInst> Instructions(SnippetInstructions, MCInstBuilder(X86::NOOP)); FunctionFiller Sink(*MF, {X86::EAX}); @@ -98,11 +101,12 @@ TEST_F(X86SnippetRepetitorTest, Loop) { HasOpcode(X86::NOOP), HasOpcode(X86::NOOP), HasOpcode(X86::NOOP), HasOpcode(X86::ADD64ri8), HasOpcode(X86::JCC_1))); - EXPECT_THAT(LoopBlock.liveins(), - UnorderedElementsAre( - LiveReg(X86::EAX), - LiveReg(State.getExegesisTarget().getLoopCounterRegister( - State.getTargetMachine().getTargetTriple())))); + EXPECT_THAT( + LoopBlock.liveins(), + UnorderedElementsAre( + LiveReg(X86::EAX), + LiveReg(State.getExegesisTarget().getDefaultLoopCounterRegister( + State.getTargetMachine().getTargetTriple())))); EXPECT_THAT(MF->getBlockNumbered(2)->instrs(), ElementsAre(HasOpcode(X86::RET64))); } diff --git a/llvm/utils/gn/secondary/llvm/lib/Target/Hexagon/BUILD.gn b/llvm/utils/gn/secondary/llvm/lib/Target/Hexagon/BUILD.gn index 747ca8f9c91d..b966b7484267 100644 --- a/llvm/utils/gn/secondary/llvm/lib/Target/Hexagon/BUILD.gn +++ b/llvm/utils/gn/secondary/llvm/lib/Target/Hexagon/BUILD.gn @@ -75,7 +75,6 @@ static_library("LLVMHexagonCodeGen") { "HexagonOptAddrMode.cpp", "HexagonOptimizeSZextends.cpp", "HexagonPeephole.cpp", - "HexagonPostIncOpt.cpp", "HexagonRDFOpt.cpp", "HexagonRegisterInfo.cpp", "HexagonSelectionDAGInfo.cpp", diff --git a/mlir/include/mlir/Dialect/SparseTensor/IR/Enums.h b/mlir/include/mlir/Dialect/SparseTensor/IR/Enums.h index cc134e7d953e..5563cb907e93 100644 --- a/mlir/include/mlir/Dialect/SparseTensor/IR/Enums.h +++ b/mlir/include/mlir/Dialect/SparseTensor/IR/Enums.h @@ -333,16 +333,28 @@ public: return lvlBits & static_cast<uint64_t>(p); } + /// Check if the `LevelType` is considered to be sparse. + constexpr bool hasSparseSemantic() const { + return isa<LevelFormat::Compressed, LevelFormat::Singleton, + LevelFormat::LooseCompressed, LevelFormat::NOutOfM>(); + } + + /// Check if the `LevelType` is considered to be dense-like. + constexpr bool hasDenseSemantic() const { + return isa<LevelFormat::Dense, LevelFormat::Batch>(); + } + /// Check if the `LevelType` needs positions array. constexpr bool isWithPosLT() const { - return isa<LevelFormat::Compressed>() || - isa<LevelFormat::LooseCompressed>(); + assert(!isa<LevelFormat::Undef>()); + return isa<LevelFormat::Compressed, LevelFormat::LooseCompressed>(); } /// Check if the `LevelType` needs coordinates array. constexpr bool isWithCrdLT() const { + assert(!isa<LevelFormat::Undef>()); // All sparse levels has coordinate array. - return !isa<LevelFormat::Dense>(); + return hasSparseSemantic(); } std::string toMLIRString() const { diff --git a/mlir/include/mlir/Dialect/SparseTensor/IR/SparseTensorAttrDefs.td b/mlir/include/mlir/Dialect/SparseTensor/IR/SparseTensorAttrDefs.td index ca98665256be..5d1db2323f95 100644 --- a/mlir/include/mlir/Dialect/SparseTensor/IR/SparseTensorAttrDefs.td +++ b/mlir/include/mlir/Dialect/SparseTensor/IR/SparseTensorAttrDefs.td @@ -374,6 +374,8 @@ def SparseTensorEncodingAttr : SparseTensor_Attr<"SparseTensorEncoding", /// is non-null (since no fixed result is valid for every dense-tensor). ::mlir::sparse_tensor::Level getLvlRank() const; + uint64_t getBatchLvlRank() const; + // // lvlTypes methods. // diff --git a/mlir/include/mlir/Dialect/SparseTensor/IR/SparseTensorStorageLayout.h b/mlir/include/mlir/Dialect/SparseTensor/IR/SparseTensorStorageLayout.h index 27dc39609cda..ce34ae43d1c1 100644 --- a/mlir/include/mlir/Dialect/SparseTensor/IR/SparseTensorStorageLayout.h +++ b/mlir/include/mlir/Dialect/SparseTensor/IR/SparseTensorStorageLayout.h @@ -30,15 +30,15 @@ namespace sparse_tensor { /// ; if dense: /// <nothing> /// ; if compressed: -/// memref<? x pos> positions ; positions for level l -/// memref<? x crd> coordinates ; coordinates for level l -/// ; if loose-compressed: -/// memref<? x pos> positions ; lo/hi position pairs for level l -/// memref<? x crd> coordinates ; coordinates for level l +/// memref<[batch] x ? x pos> positions ; positions for level l +/// memref<[batch] x ? x crd> coordinates ; coordinates for level l +/// ; if loose-[batch] x compressed: +/// memref<[batch] x ? x pos> positions ; lo/hi pos pairs for level l +/// memref<[batch] x ? x crd> coordinates ; coordinates for level l /// ; if singleton/2-out-of-4: -/// memref<? x crd> coordinates ; coordinates for level l +/// memref<[batch] x ? x crd> coordinates ; coordinates for level l /// -/// memref<? x eltType> values ; values +/// memref<[batch] x ? x eltType> values ; values /// /// struct sparse_tensor.storage_specifier { /// array<rank x int> lvlSizes ; sizes/cardinalities for each level diff --git a/mlir/include/mlir/Dialect/SparseTensor/IR/SparseTensorType.h b/mlir/include/mlir/Dialect/SparseTensor/IR/SparseTensorType.h index 1a090ddb782f..c93a4fcd922c 100644 --- a/mlir/include/mlir/Dialect/SparseTensor/IR/SparseTensorType.h +++ b/mlir/include/mlir/Dialect/SparseTensor/IR/SparseTensorType.h @@ -253,6 +253,14 @@ public: CrdTransDirectionKind::dim2lvl); } + /// Returns the Level-shape. + SmallVector<Size> getBatchLvlShape() const { + auto lvlShape = getEncoding().tranlateShape(getDimShape(), + CrdTransDirectionKind::dim2lvl); + lvlShape.truncate(getEncoding().getBatchLvlRank()); + return lvlShape; + } + /// Returns the type with an identity mapping. RankedTensorType getDemappedType() const { return RankedTensorType::get(getLvlShape(), getElementType(), diff --git a/mlir/include/mlir/Dialect/SparseTensor/Utils/Merger.h b/mlir/include/mlir/Dialect/SparseTensor/Utils/Merger.h index 490ef3071af1..7f9820df984b 100644 --- a/mlir/include/mlir/Dialect/SparseTensor/Utils/Merger.h +++ b/mlir/include/mlir/Dialect/SparseTensor/Utils/Merger.h @@ -509,8 +509,7 @@ public: bool isSparseLvlWithNonTrivialIdxExp(TensorLoopId b) const { if (isLvlWithNonTrivialIdxExp(b)) { auto lt = getLoopDependentLevelType(b); - return isCompressedLT(lt) || isSingletonLT(lt) || - isLooseCompressedLT(lt) || isNOutOfMLT(lt); + return lt.hasSparseSemantic(); } return false; } diff --git a/mlir/include/mlir/IR/Value.h b/mlir/include/mlir/IR/Value.h index fff3b87faff6..a74d0faa1dfc 100644 --- a/mlir/include/mlir/IR/Value.h +++ b/mlir/include/mlir/IR/Value.h @@ -90,7 +90,7 @@ protected: /// class has value-type semantics and is just a simple wrapper around a /// ValueImpl that is either owner by a block(in the case of a BlockArgument) or /// an Operation(in the case of an OpResult). -/// As most IR construct, this isn't const-correct, but we keep method +/// As most IR constructs, this isn't const-correct, but we keep method /// consistent and as such method that immediately modify this Value aren't /// marked `const` (include modifying the Value use-list). class Value { diff --git a/mlir/lib/Dialect/SparseTensor/IR/SparseTensorDialect.cpp b/mlir/lib/Dialect/SparseTensor/IR/SparseTensorDialect.cpp index fd0ed26fbde0..69c3413f35ea 100644 --- a/mlir/lib/Dialect/SparseTensor/IR/SparseTensorDialect.cpp +++ b/mlir/lib/Dialect/SparseTensor/IR/SparseTensorDialect.cpp @@ -126,13 +126,16 @@ void sparse_tensor::foreachFieldAndTypeInSparseTensor( const Type posType = stt.getPosType(); const Type eltType = stt.getElementType(); + SmallVector<int64_t> memrefShape = stt.getBatchLvlShape(); + memrefShape.push_back(ShapedType::kDynamic); + const Type specType = StorageSpecifierType::get(stt.getEncoding()); - // memref<? x pos> positions - const Type posMemType = MemRefType::get({ShapedType::kDynamic}, posType); - // memref<? x crd> coordinates - const Type crdMemType = MemRefType::get({ShapedType::kDynamic}, crdType); - // memref<? x eltType> values - const Type valMemType = MemRefType::get({ShapedType::kDynamic}, eltType); + // memref<[batch] x ? x pos> positions + const Type posMemType = MemRefType::get(memrefShape, posType); + // memref<[batch] x ? x crd> coordinates + const Type crdMemType = MemRefType::get(memrefShape, crdType); + // memref<[batch] x ? x eltType> values + const Type valMemType = MemRefType::get(memrefShape, eltType); StorageLayout(stt).foreachField([specType, posMemType, crdMemType, valMemType, callback](FieldIndex fieldIdx, @@ -336,6 +339,12 @@ SparseTensorEncodingAttr SparseTensorEncodingAttr::withoutDimSlices() const { return withDimSlices(ArrayRef<SparseTensorDimSliceAttr>{}); } +uint64_t SparseTensorEncodingAttr::getBatchLvlRank() const { + ArrayRef<LevelType> lvlTypes = getLvlTypes(); + auto lastBatch = std::find_if(lvlTypes.rbegin(), lvlTypes.rend(), isBatchLT); + return std::distance(lastBatch, lvlTypes.rend()); +} + bool SparseTensorEncodingAttr::isAllDense() const { return !getImpl() || llvm::all_of(getLvlTypes(), isDenseLT); } diff --git a/mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorCodegen.cpp b/mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorCodegen.cpp index 0ccb11f3a6b8..d5eec4ae67e7 100644 --- a/mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorCodegen.cpp +++ b/mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorCodegen.cpp @@ -1293,7 +1293,7 @@ struct SparseAssembleOpConverter : public OpConversionPattern<AssembleOp> { Value tensor = fKind == SparseTensorFieldKind::ValMemRef ? op.getValues() : op.getLevels()[fIdx]; - + // TODO: handle batch. TypedValue<BaseMemRefType> mem = genToMemref(rewriter, loc, tensor); if (mem.getType().getRank() > 1) { // Flattens the buffer to rank 1. @@ -1322,9 +1322,8 @@ struct SparseAssembleOpConverter : public OpConversionPattern<AssembleOp> { for (Level lvl = 0, lvlRank = stt.getLvlRank(); lvl < lvlRank; lvl++) { assert(!ShapedType::isDynamic(stt.getDimShape()[lvl])); - // FIXME: dim/lvl confusion! // Sets up the level size. - auto lvlSize = constantIndex(rewriter, loc, stt.getDimShape()[lvl]); + auto lvlSize = constantIndex(rewriter, loc, stt.getLvlShape()[lvl]); desc.setLvlSize(rewriter, loc, lvl, lvlSize); // We use a single AOS array to store the trailing COO, so there is only // one memory size to set for the entire COO section. diff --git a/mlir/lib/Dialect/SparseTensor/Utils/Merger.cpp b/mlir/lib/Dialect/SparseTensor/Utils/Merger.cpp index 731cd79a1e3b..72b722c69ae3 100644 --- a/mlir/lib/Dialect/SparseTensor/Utils/Merger.cpp +++ b/mlir/lib/Dialect/SparseTensor/Utils/Merger.cpp @@ -476,7 +476,7 @@ BitVector Merger::simplifyCond(LatSetId s0, LatPointId p0) { // Starts resetting from a dense level, so that the first bit (if kept) // is not undefined level-type. for (unsigned b = 0; b < be; b++) { - if (simple[b] && isDenseLT(getLvlType(TensorLoopId{b}))) { + if (simple[b] && getLvlType(TensorLoopId{b}).hasDenseSemantic()) { offset = be - b - 1; // relative to the end break; } @@ -489,8 +489,7 @@ BitVector Merger::simplifyCond(LatSetId s0, LatPointId p0) { // Slice on dense level has `locate` property as well, and can be optimized. if (simple[b] && !isSparseLvlWithNonTrivialIdxExp(b)) { const auto lt = getLvlType(b); - if (!isCompressedLT(lt) && !isSingletonLT(lt) && - !isLooseCompressedLT(lt) && !isNOutOfMLT(lt)) { + if (!lt.hasSparseSemantic()) { if (reset) simple.reset(b); reset = true; @@ -670,8 +669,7 @@ bool Merger::isSingleCondition(TensorId t, ExprId e) const { bool Merger::hasAnySparse(const BitVector &bits) const { for (TensorLoopId b : bits.set_bits()) { const auto lt = getLvlType(b); - if (isCompressedLT(lt) || isSingletonLT(lt) || isLooseCompressedLT(lt) || - isNOutOfMLT(lt)) + if (lt.hasSparseSemantic()) return true; } return hasSparseIdxReduction(bits); diff --git a/mlir/test/Dialect/SparseTensor/codegen.mlir b/mlir/test/Dialect/SparseTensor/codegen.mlir index c1a976c84fec..b63762485c96 100644 --- a/mlir/test/Dialect/SparseTensor/codegen.mlir +++ b/mlir/test/Dialect/SparseTensor/codegen.mlir @@ -34,6 +34,10 @@ map = (d0, d1) -> (d1 : dense, d0 : compressed) }> +#BCSR = #sparse_tensor.encoding<{ + map = (d0, d1, d2, d3) -> (d0: batch, d1: batch, d2 : dense, d3 : compressed) +}> + #DCSR = #sparse_tensor.encoding<{ map = (d0, d1) -> (d0 : compressed, d1 : compressed), crdWidth = 64, @@ -182,6 +186,36 @@ func.func @sparse_csr(%arg0: tensor<?x?xf64, #CSR>) { return } +// CHECK-LABEL: func @sparse_bcsr_0( +// CHECK-SAME: %[[A1:.*0]]: memref<?x2x?xindex>, +// CHECK-SAME: %[[A2:.*1]]: memref<?x2x?xindex>, +// CHECK-SAME: %[[A3:.*]]: memref<?x2x?xf64>, +// CHECK-SAME: %[[A4:.*]]: !sparse_tensor.storage_specifier +// CHECK: return +func.func @sparse_bcsr_0(%arg0: tensor<?x2x?x?xf64, #BCSR>) { + return +} + +// CHECK-LABEL: func @sparse_bcsr_1( +// CHECK-SAME: %[[A1:.*0]]: memref<?x?x?xindex>, +// CHECK-SAME: %[[A2:.*1]]: memref<?x?x?xindex>, +// CHECK-SAME: %[[A3:.*]]: memref<?x?x?xf64>, +// CHECK-SAME: %[[A4:.*]]: !sparse_tensor.storage_specifier +// CHECK: return +func.func @sparse_bcsr_1(%arg0: tensor<?x?x?x?xf64, #BCSR>) { + return +} + +// CHECK-LABEL: func @sparse_bcsr_2( +// CHECK-SAME: %[[A1:.*0]]: memref<18x6x?xindex>, +// CHECK-SAME: %[[A2:.*1]]: memref<18x6x?xindex>, +// CHECK-SAME: %[[A3:.*]]: memref<18x6x?xf64>, +// CHECK-SAME: %[[A4:.*]]: !sparse_tensor.storage_specifier +// CHECK: return +func.func @sparse_bcsr_2(%arg0: tensor<18x6x4x2xf64, #BCSR>) { + return +} + // CHECK-LABEL: func @sparse_dcsr( // CHECK-SAME: %[[A0:.*0]]: memref<?xi32>, // CHECK-SAME: %[[A1:.*1]]: memref<?xi64>, diff --git a/mlir/unittests/Dialect/SparseTensor/MergerTest.cpp b/mlir/unittests/Dialect/SparseTensor/MergerTest.cpp index 62a19c084cac..943e7d5c120b 100644 --- a/mlir/unittests/Dialect/SparseTensor/MergerTest.cpp +++ b/mlir/unittests/Dialect/SparseTensor/MergerTest.cpp @@ -120,7 +120,8 @@ static Match synZeroMatch() { return Match(); } FOREVERY_BINOP(IMPL_BINOP_PATTERN) #undef IMPL_BINOP_PATTERN -class MergerTestBase : public ::testing::Test { +// Parameterize LevelFormat to test both Dense and Batch LevelFormat. +class MergerTestBase : public ::testing::TestWithParam<LevelFormat> { protected: MergerTestBase(unsigned numTensors, unsigned numLoops) : merger(numTensors, numLoops, /*maxRank=*/numLoops) { @@ -317,10 +318,14 @@ protected: // Tensor 1: sparse input vector. merger.setLevelAndType(tid(1), lid(0), 0, LevelFormat::Compressed); // Tensor 2: dense output vector. - merger.setLevelAndType(tid(2), lid(0), 0, LevelFormat::Dense); + merger.setLevelAndType(tid(2), lid(0), 0, GetParam()); } }; +INSTANTIATE_TEST_SUITE_P(Test3T1L, MergerTest3T1L, + ::testing::Values(LevelFormat::Dense, + LevelFormat::Batch)); + /// Four tensors (three inputs, one output); and a single loop. class MergerTest4T1L : public MergerTestBase { protected: @@ -333,10 +338,14 @@ protected: // Tensor 2: sparse input vector merger.setLevelAndType(tid(2), lid(0), 0, LevelFormat::Compressed); // Tensor 3: dense output vector - merger.setLevelAndType(tid(3), lid(0), 0, LevelFormat::Dense); + merger.setLevelAndType(tid(3), lid(0), 0, GetParam()); } }; +INSTANTIATE_TEST_SUITE_P(Test4T1L, MergerTest4T1L, + ::testing::Values(LevelFormat::Dense, + LevelFormat::Batch)); + /// /// Tests with both sparse and dense input. /// @@ -349,12 +358,16 @@ protected: // Tensor 0: sparse input vector. merger.setLevelAndType(tid(0), lid(0), 0, LevelFormat::Compressed); // Tensor 1: dense input vector. - merger.setLevelAndType(tid(1), lid(0), 0, LevelFormat::Dense); + merger.setLevelAndType(tid(1), lid(0), 0, GetParam()); // Tensor 2: dense output vector. - merger.setLevelAndType(tid(2), lid(0), 0, LevelFormat::Dense); + merger.setLevelAndType(tid(2), lid(0), 0, GetParam()); } }; +INSTANTIATE_TEST_SUITE_P(Test3T1LD, MergerTest3T1LD, + ::testing::Values(LevelFormat::Dense, + LevelFormat::Batch)); + /// /// Tests with both undef and dense input. /// @@ -367,14 +380,18 @@ protected: // Tensor 0: undef input vector. merger.setLevelAndType(tid(0), lid(0), 0, LevelFormat::Undef); // Tensor 1: dense input vector. - merger.setLevelAndType(tid(1), lid(0), 0, LevelFormat::Dense); + merger.setLevelAndType(tid(1), lid(0), 0, GetParam()); // Tensor 2: undef input vector. merger.setLevelAndType(tid(2), lid(0), 0, LevelFormat::Undef); // Tensor 3: dense output vector. - merger.setLevelAndType(tid(3), lid(0), 0, LevelFormat::Dense); + merger.setLevelAndType(tid(3), lid(0), 0, GetParam()); } }; +INSTANTIATE_TEST_SUITE_P(Test4T1LU, MergerTest4T1LU, + ::testing::Values(LevelFormat::Dense, + LevelFormat::Batch)); + /// /// Tests with operation on sparse output. /// @@ -395,6 +412,11 @@ protected: } }; +// This testsuite does not use any dense-like format, just one of {Dense, Batch} +// is enough. +INSTANTIATE_TEST_SUITE_P(Test3T1LSo, MergerTest3T1LSo, + ::testing::Values(LevelFormat::Dense)); + } // namespace /// Vector multiplication (conjunction) of 3 vectors, i.e.; @@ -409,7 +431,7 @@ protected: /// lat( i_01_D / (tensor_0 * tensor_1 * tensor2) ) /// } #define IMPL_MERGER_TEST_CONJ_CONJ_UNDEF(CONJ1, CONJ2) \ - TEST_F(MergerTest4T1LU, vector_##CONJ1##_##CONJ2) { \ + TEST_P(MergerTest4T1LU, vector_##CONJ1##_##CONJ2) { \ const auto em = CONJ1##Expr(tensor(0), tensor(1)); \ const auto e = CONJ2##Expr(em, tensor(2)); \ const auto l0 = lid(0); \ @@ -443,7 +465,7 @@ FOREVERY_PAIR_OF_COMMON_CONJ_CONJ_BINOP(IMPL_MERGER_TEST_CONJ_CONJ_UNDEF) /// lat( i_03_U / (tensor_0 * tensor_1 * output_tensor2) ) /// } #define IMPL_MERGER_TEST_CONJ_CONJ_SPARSE_OUT(CONJ1, CONJ2) \ - TEST_F(MergerTest3T1LSo, vector_##CONJ1##_##CONJ2) { \ + TEST_P(MergerTest3T1LSo, vector_##CONJ1##_##CONJ2) { \ const auto em = CONJ1##Expr(tensor(0), tensor(1)); \ const auto e = CONJ2##Expr(em, tensor(2)); \ const auto l0 = lid(0); \ @@ -482,7 +504,7 @@ FOREVERY_PAIR_OF_COMMON_CONJ_CONJ_BINOP(IMPL_MERGER_TEST_CONJ_CONJ_SPARSE_OUT) /// lat( i_01 / tensor_1 ) /// } #define IMPL_MERGER_TEST_DISJ(OP, UNUSED) \ - TEST_F(MergerTest3T1L, vector_##OP) { \ + TEST_P(MergerTest3T1L, vector_##OP) { \ const auto e = OP##Expr(tensor(0), tensor(1)); \ const auto l0 = lid(0); \ const auto t0 = tid(0); \ @@ -514,7 +536,7 @@ FOREVERY_COMMON_DISJ_BINOP(IMPL_MERGER_TEST_DISJ) /// lat( i_00 i_01 / (tensor_0 * tensor_1) ) /// } #define IMPL_MERGER_TEST_CONJ(OP, UNUSED) \ - TEST_F(MergerTest3T1L, vector_##OP) { \ + TEST_P(MergerTest3T1L, vector_##OP) { \ const auto e = OP##Expr(tensor(0), tensor(1)); \ const auto l0 = lid(0); \ const auto t0 = tid(0); \ @@ -544,7 +566,7 @@ FOREVERY_COMMON_CONJ_BINOP(IMPL_MERGER_TEST_CONJ) /// lat( i_02 / tensor_2 ) /// } #define IMPL_MERGER_TEST_CONJ_DISJ(CONJ, DISJ) \ - TEST_F(MergerTest4T1L, vector_##CONJ##_##DISJ) { \ + TEST_P(MergerTest4T1L, vector_##CONJ##_##DISJ) { \ const auto em = CONJ##Expr(tensor(0), tensor(1)); \ const auto e = DISJ##Expr(em, tensor(2)); \ const auto l0 = lid(0); \ @@ -587,7 +609,7 @@ FOREVERY_PAIR_OF_COMMON_CONJ_DISJ_BINOP(IMPL_MERGER_TEST_CONJ_DISJ) /// lat( i_00 / tensor_0 ) /// } #define IMPL_MERGER_TEST_DISJ_DISJ(DISJ1, DISJ2) \ - TEST_F(MergerTest4T1L, Vector_##DISJ1##_##DISJ2) { \ + TEST_P(MergerTest4T1L, Vector_##DISJ1##_##DISJ2) { \ const auto em = DISJ1##Expr(tensor(0), tensor(1)); \ const auto e = DISJ2##Expr(em, tensor(2)); \ const auto l0 = lid(0); \ @@ -636,7 +658,7 @@ FOREVERY_PAIR_OF_COMMON_DISJ_DISJ_BINOP(IMPL_MERGER_TEST_DISJ_DISJ) /// lat( i_00 i_01 i_02 / tensor_0 * tensor_1 * tensor_2 ) /// } #define IMPL_MERGER_TEST_CONJ_CONJ(CONJ1, CONJ2) \ - TEST_F(MergerTest4T1L, vector_##CONJ1##_##CONJ2) { \ + TEST_P(MergerTest4T1L, vector_##CONJ1##_##CONJ2) { \ const auto em = CONJ1##Expr(tensor(0), tensor(1)); \ const auto e = CONJ2##Expr(em, tensor(2)); \ const auto l0 = lid(0); \ @@ -675,7 +697,7 @@ FOREVERY_PAIR_OF_COMMON_CONJ_CONJ_BINOP(IMPL_MERGER_TEST_CONJ_CONJ) /// lat( i_00 / sparse_tensor_0 ) should be opted out as it only has dense diff /// with lat( i_00 i_01 / (sparse_tensor_0 + dense_tensor_1) ). #define IMPL_MERGER_TEST_OPTIMIZED_DISJ(OP, UNUSED) \ - TEST_F(MergerTest3T1LD, vector_opted_##OP) { \ + TEST_P(MergerTest3T1LD, vector_opted_##OP) { \ const auto e = OP##Expr(tensor(0), tensor(1)); \ const auto l0 = lid(0); \ const auto t0 = tid(0); \ @@ -711,7 +733,7 @@ FOREVERY_COMMON_DISJ_BINOP(IMPL_MERGER_TEST_OPTIMIZED_DISJ) /// } /// since i_01 is a dense dimension. #define IMPL_MERGER_TEST_OPTIMIZED_CONJ(OP, UNUSED) \ - TEST_F(MergerTest3T1LD, vector_opted_##OP) { \ + TEST_P(MergerTest3T1LD, vector_opted_##OP) { \ const auto e = OP##Expr(tensor(0), tensor(1)); \ const auto l0 = lid(0); \ const auto t0 = tid(0); \ @@ -746,7 +768,7 @@ FOREVERY_COMMON_CONJ_BINOP(IMPL_MERGER_TEST_OPTIMIZED_CONJ) /// lat( i_00 / tensor_0 cmp 0 ) /// lat( i_01 / 0 cmp tensor_1 ) /// } -TEST_F(MergerTest3T1L, vector_cmp) { +TEST_P(MergerTest3T1L, vector_cmp) { const auto e = cmpiExpr(tensor(0), tensor(1)); const auto l0 = lid(0); const auto t0 = tid(0); @@ -784,7 +806,7 @@ TEST_F(MergerTest3T1L, vector_cmp) { /// /// lat( i_00 / sparse_tensor_0 ) should be opted out as it only has dense diff /// with lat( i_00 i_01 / (sparse_tensor_0 cmp dense_tensor_1) ). -TEST_F(MergerTest3T1LD, vector_cmp) { +TEST_P(MergerTest3T1LD, vector_cmp) { const auto e = cmpiExpr(tensor(0), tensor(1)); const auto l0 = lid(0); const auto t0 = tid(0); diff --git a/openmp/runtime/src/kmp_runtime.cpp b/openmp/runtime/src/kmp_runtime.cpp index fc5e8405a415..7edb0b440acc 100644 --- a/openmp/runtime/src/kmp_runtime.cpp +++ b/openmp/runtime/src/kmp_runtime.cpp @@ -5708,9 +5708,8 @@ void __kmp_free_team(kmp_root_t *root, } #endif // first check if thread is sleeping - kmp_flag_64<> fl(&th->th.th_bar[bs_forkjoin_barrier].bb.b_go, th); - if (fl.is_sleeping()) - fl.resume(__kmp_gtid_from_thread(th)); + if (th->th.th_sleep_loc) + __kmp_null_resume_wrapper(th); KMP_CPU_PAUSE(); } } diff --git a/openmp/runtime/test/barrier/llvm-issue-80664.c b/openmp/runtime/test/barrier/llvm-issue-80664.c new file mode 100644 index 000000000000..79aa228afa6b --- /dev/null +++ b/openmp/runtime/test/barrier/llvm-issue-80664.c @@ -0,0 +1,37 @@ +// RUN: %libomp-compile +// RUN: env OMP_WAIT_POLICY=passive \ +// RUN: KMP_FORKJOIN_BARRIER_PATTERN='linear,linear' %libomp-run +// RUN: env OMP_WAIT_POLICY=passive \ +// RUN: KMP_FORKJOIN_BARRIER_PATTERN='tree,tree' %libomp-run +// RUN: env OMP_WAIT_POLICY=passive \ +// RUN: KMP_FORKJOIN_BARRIER_PATTERN='hyper,hyper' %libomp-run +// RUN: env OMP_WAIT_POLICY=passive \ +// RUN: KMP_FORKJOIN_BARRIER_PATTERN='dist,dist' %libomp-run +// +// LLVM ISSUE 80664: https://github.com/llvm/llvm-project/issues/80664 +// +// Distributed barrier + OMP_WAIT_POLICY=passive hangs in library termination +// Reason: the resume logic in __kmp_free_team() was faulty and, when checking +// for sleep status, didn't look at correct location for distributed barrier. + +#include <stdio.h> +#include <stdlib.h> + +int a = 0; + +void test_omp_barrier() { +#pragma omp parallel + { +#pragma omp task + { +#pragma omp atomic + a++; + } + } +} + +int main() { + test_omp_barrier(); + printf("a = %d\n", a); + return EXIT_SUCCESS; +} |