diff --git a/clang-tools-extra/clang-tidy/bugprone/SuspiciousSemicolonCheck.cpp b/clang-tools-extra/clang-tidy/bugprone/SuspiciousSemicolonCheck.cpp index 9b34f5ab55a7fc..3874a52e7e20a2 100644 --- a/clang-tools-extra/clang-tidy/bugprone/SuspiciousSemicolonCheck.cpp +++ b/clang-tools-extra/clang-tidy/bugprone/SuspiciousSemicolonCheck.cpp @@ -54,10 +54,10 @@ void SuspiciousSemicolonCheck::check(const MatchFinder::MatchResult &Result) { SourceLocation LocEnd = Semicolon->getEndLoc(); FileID FID = SM.getFileID(LocEnd); - const llvm::MemoryBuffer *Buffer = SM.getBuffer(FID, LocEnd); + llvm::MemoryBufferRef Buffer = SM.getBufferOrFake(FID, LocEnd); Lexer Lexer(SM.getLocForStartOfFile(FID), Ctxt.getLangOpts(), - Buffer->getBufferStart(), SM.getCharacterData(LocEnd) + 1, - Buffer->getBufferEnd()); + Buffer.getBufferStart(), SM.getCharacterData(LocEnd) + 1, + Buffer.getBufferEnd()); if (Lexer.LexFromRawLexer(Token)) return; diff --git a/clang-tools-extra/clang-tidy/misc/RedundantExpressionCheck.cpp b/clang-tools-extra/clang-tidy/misc/RedundantExpressionCheck.cpp index b84e4d525d8cf2..5d3d3c46edb4a6 100644 --- a/clang-tools-extra/clang-tidy/misc/RedundantExpressionCheck.cpp +++ b/clang-tools-extra/clang-tidy/misc/RedundantExpressionCheck.cpp @@ -794,14 +794,14 @@ static bool areExprsFromDifferentMacros(const Expr *LhsExpr, SM.getDecomposedLoc(SM.getExpansionLoc(Lsr.getBegin())); std::pair RsrLocInfo = SM.getDecomposedLoc(SM.getExpansionLoc(Rsr.getBegin())); - const llvm::MemoryBuffer *MB = SM.getBuffer(LsrLocInfo.first); + llvm::MemoryBufferRef MB = SM.getBufferOrFake(LsrLocInfo.first); - const char *LTokenPos = MB->getBufferStart() + LsrLocInfo.second; - const char *RTokenPos = MB->getBufferStart() + RsrLocInfo.second; + const char *LTokenPos = MB.getBufferStart() + LsrLocInfo.second; + const char *RTokenPos = MB.getBufferStart() + RsrLocInfo.second; Lexer LRawLex(SM.getLocForStartOfFile(LsrLocInfo.first), LO, - MB->getBufferStart(), LTokenPos, MB->getBufferEnd()); + MB.getBufferStart(), LTokenPos, MB.getBufferEnd()); Lexer RRawLex(SM.getLocForStartOfFile(RsrLocInfo.first), LO, - MB->getBufferStart(), RTokenPos, MB->getBufferEnd()); + MB.getBufferStart(), RTokenPos, MB.getBufferEnd()); Token LTok, RTok; do { // Compare the expressions token-by-token. diff --git a/clang-tools-extra/clang-tidy/misc/StaticAssertCheck.cpp b/clang-tools-extra/clang-tidy/misc/StaticAssertCheck.cpp index 53cebc32bb38ad..3e3f8dbf02ff55 100644 --- a/clang-tools-extra/clang-tidy/misc/StaticAssertCheck.cpp +++ b/clang-tools-extra/clang-tidy/misc/StaticAssertCheck.cpp @@ -140,7 +140,8 @@ SourceLocation StaticAssertCheck::getLastParenLoc(const ASTContext *ASTCtx, const LangOptions &Opts = ASTCtx->getLangOpts(); const SourceManager &SM = ASTCtx->getSourceManager(); - const llvm::MemoryBuffer *Buffer = SM.getBuffer(SM.getFileID(AssertLoc)); + llvm::Optional Buffer = + SM.getBufferOrNone(SM.getFileID(AssertLoc)); if (!Buffer) return SourceLocation(); diff --git a/clang-tools-extra/clang-tidy/readability/MagicNumbersCheck.cpp b/clang-tools-extra/clang-tidy/readability/MagicNumbersCheck.cpp index 1cecfa31da3428..18165a619adf45 100644 --- a/clang-tools-extra/clang-tidy/readability/MagicNumbersCheck.cpp +++ b/clang-tools-extra/clang-tidy/readability/MagicNumbersCheck.cpp @@ -209,7 +209,7 @@ bool MagicNumbersCheck::isSyntheticValue(const SourceManager *SourceManager, return false; const StringRef BufferIdentifier = - SourceManager->getBuffer(FileOffset.first)->getBufferIdentifier(); + SourceManager->getBufferOrFake(FileOffset.first).getBufferIdentifier(); return BufferIdentifier.empty(); } diff --git a/clang-tools-extra/clangd/ParsedAST.cpp b/clang-tools-extra/clangd/ParsedAST.cpp index 23d86221e74fee..c8096017b50ab0 100644 --- a/clang-tools-extra/clangd/ParsedAST.cpp +++ b/clang-tools-extra/clangd/ParsedAST.cpp @@ -164,7 +164,7 @@ class ReplayPreamble : private PPCallbacks { SrcMgr::CharacteristicKind Kind, FileID PrevFID) override { // It'd be nice if there was a better way to identify built-in headers... if (Reason == FileChangeReason::ExitFile && - SM.getBuffer(PrevFID)->getBufferIdentifier() == "") + SM.getBufferOrFake(PrevFID).getBufferIdentifier() == "") replay(); } diff --git a/clang-tools-extra/clangd/SemanticHighlighting.cpp b/clang-tools-extra/clangd/SemanticHighlighting.cpp index 4e66a9bb4e857e..a9c885c7275e9c 100644 --- a/clang-tools-extra/clangd/SemanticHighlighting.cpp +++ b/clang-tools-extra/clangd/SemanticHighlighting.cpp @@ -222,7 +222,7 @@ class HighlightingsBuilder { TokRef = TokRef.drop_front(Conflicting.size()); } const auto &SM = AST.getSourceManager(); - StringRef MainCode = SM.getBuffer(SM.getMainFileID())->getBuffer(); + StringRef MainCode = SM.getBufferOrFake(SM.getMainFileID()).getBuffer(); // Merge token stream with "inactive line" markers. std::vector WithInactiveLines; diff --git a/clang-tools-extra/clangd/SourceCode.cpp b/clang-tools-extra/clangd/SourceCode.cpp index c6279177eba99b..bcd5758af44118 100644 --- a/clang-tools-extra/clangd/SourceCode.cpp +++ b/clang-tools-extra/clangd/SourceCode.cpp @@ -445,9 +445,8 @@ llvm::Optional toHalfOpenFileRange(const SourceManager &SM, llvm::StringRef toSourceCode(const SourceManager &SM, SourceRange R) { assert(isValidFileRange(SM, R)); - bool Invalid = false; - auto *Buf = SM.getBuffer(SM.getFileID(R.getBegin()), &Invalid); - assert(!Invalid); + auto Buf = SM.getBufferOrNone(SM.getFileID(R.getBegin())); + assert(Buf); size_t BeginOffset = SM.getFileOffset(R.getBegin()); size_t EndOffset = SM.getFileOffset(R.getEnd()); @@ -456,7 +455,7 @@ llvm::StringRef toSourceCode(const SourceManager &SM, SourceRange R) { llvm::Expected sourceLocationInMainFile(const SourceManager &SM, Position P) { - llvm::StringRef Code = SM.getBuffer(SM.getMainFileID())->getBuffer(); + llvm::StringRef Code = SM.getBufferOrFake(SM.getMainFileID()).getBuffer(); auto Offset = positionToOffset(Code, P, /*AllowColumnBeyondLineLength=*/false); if (!Offset) diff --git a/clang-tools-extra/modularize/PreprocessorTracker.cpp b/clang-tools-extra/modularize/PreprocessorTracker.cpp index f8ab2c8067c010..f472415225a172 100644 --- a/clang-tools-extra/modularize/PreprocessorTracker.cpp +++ b/clang-tools-extra/modularize/PreprocessorTracker.cpp @@ -312,10 +312,10 @@ static std::string getSourceString(clang::Preprocessor &PP, // Retrieve source line from file image given a location. static std::string getSourceLine(clang::Preprocessor &PP, clang::SourceLocation Loc) { - const llvm::MemoryBuffer *MemBuffer = - PP.getSourceManager().getBuffer(PP.getSourceManager().getFileID(Loc)); - const char *Buffer = MemBuffer->getBufferStart(); - const char *BufferEnd = MemBuffer->getBufferEnd(); + llvm::MemoryBufferRef MemBuffer = PP.getSourceManager().getBufferOrFake( + PP.getSourceManager().getFileID(Loc)); + const char *Buffer = MemBuffer.getBufferStart(); + const char *BufferEnd = MemBuffer.getBufferEnd(); const char *BeginPtr = PP.getSourceManager().getCharacterData(Loc); const char *EndPtr = BeginPtr; while (BeginPtr > Buffer) { @@ -338,9 +338,10 @@ static std::string getSourceLine(clang::Preprocessor &PP, // Retrieve source line from file image given a file ID and line number. static std::string getSourceLine(clang::Preprocessor &PP, clang::FileID FileID, int Line) { - const llvm::MemoryBuffer *MemBuffer = PP.getSourceManager().getBuffer(FileID); - const char *Buffer = MemBuffer->getBufferStart(); - const char *BufferEnd = MemBuffer->getBufferEnd(); + llvm::MemoryBufferRef MemBuffer = + PP.getSourceManager().getBufferOrFake(FileID); + const char *Buffer = MemBuffer.getBufferStart(); + const char *BufferEnd = MemBuffer.getBufferEnd(); const char *BeginPtr = Buffer; const char *EndPtr = BufferEnd; int LineCounter = 1; diff --git a/clang/include/clang/Basic/SourceManager.h b/clang/include/clang/Basic/SourceManager.h index e34fe4577aae85..2156a013e53b7e 100644 --- a/clang/include/clang/Basic/SourceManager.h +++ b/clang/include/clang/Basic/SourceManager.h @@ -980,6 +980,17 @@ class SourceManager : public RefCountedBase { Diag, getFileManager(), Loc); } + /// Return the buffer for the specified FileID. + /// + /// If there is an error opening this buffer the first time, this + /// manufactures a temporary buffer and returns it. + llvm::MemoryBufferRef + getBufferOrFake(FileID FID, SourceLocation Loc = SourceLocation()) const { + if (auto B = getBufferOrNone(FID, Loc)) + return *B; + return getFakeBufferForRecovery()->getMemBufferRef(); + } + /// Return the buffer for the specified FileID. /// /// If there is an error opening this buffer the first time, this diff --git a/clang/lib/Basic/Diagnostic.cpp b/clang/lib/Basic/Diagnostic.cpp index 661eabf9bc7cbb..ba8c2fb4731dd7 100644 --- a/clang/lib/Basic/Diagnostic.cpp +++ b/clang/lib/Basic/Diagnostic.cpp @@ -265,7 +265,8 @@ void DiagnosticsEngine::DiagStateMap::dump(SourceManager &SrcMgr, PrintedOuterHeading = true; llvm::errs() << "File " << &File << " : " << SrcMgr.getBuffer(ID)->getBufferIdentifier(); + << ">: " << SrcMgr.getBufferOrFake(ID).getBufferIdentifier(); + if (F.second.Parent) { std::pair Decomp = SrcMgr.getDecomposedIncludedLoc(ID); diff --git a/clang/lib/Basic/SourceLocation.cpp b/clang/lib/Basic/SourceLocation.cpp index c1fa406909fe4b..8cb0899ea39dd3 100644 --- a/clang/lib/Basic/SourceLocation.cpp +++ b/clang/lib/Basic/SourceLocation.cpp @@ -245,7 +245,7 @@ const char *FullSourceLoc::getCharacterData(bool *Invalid) const { StringRef FullSourceLoc::getBufferData(bool *Invalid) const { assert(isValid()); - return SrcMgr->getBuffer(SrcMgr->getFileID(*this), Invalid)->getBuffer(); + return SrcMgr->getBufferData(SrcMgr->getFileID(*this), Invalid); } std::pair FullSourceLoc::getDecomposedLoc() const { diff --git a/clang/lib/Basic/SourceManager.cpp b/clang/lib/Basic/SourceManager.cpp index 1e1915198cd046..61e186e6aa4893 100644 --- a/clang/lib/Basic/SourceManager.cpp +++ b/clang/lib/Basic/SourceManager.cpp @@ -1228,12 +1228,11 @@ const char *SourceManager::getCharacterData(SourceLocation SL, /// this is significantly cheaper to compute than the line number. unsigned SourceManager::getColumnNumber(FileID FID, unsigned FilePos, bool *Invalid) const { - bool MyInvalid = false; - const llvm::MemoryBuffer *MemBuf = getBuffer(FID, &MyInvalid); + llvm::Optional MemBuf = getBufferOrNone(FID); if (Invalid) - *Invalid = MyInvalid; + *Invalid = !MemBuf; - if (MyInvalid) + if (!MemBuf) return 1; // It is okay to request a position just past the end of the buffer. @@ -1509,7 +1508,10 @@ StringRef SourceManager::getBufferName(SourceLocation Loc, bool *Invalid) const { if (isInvalid(Loc, Invalid)) return ""; - return getBuffer(getFileID(Loc), Invalid)->getBufferIdentifier(); + auto B = getBufferOrNone(getFileID(Loc)); + if (Invalid) + *Invalid = !B; + return B ? B->getBufferIdentifier() : ""; } /// getPresumedLoc - This method returns the "presumed" location of a @@ -2047,8 +2049,8 @@ bool SourceManager::isBeforeInTranslationUnit(SourceLocation LHS, // If we arrived here, the location is either in a built-ins buffer or // associated with global inline asm. PR5662 and PR22576 are examples. - StringRef LB = getBuffer(LOffs.first)->getBufferIdentifier(); - StringRef RB = getBuffer(ROffs.first)->getBufferIdentifier(); + StringRef LB = getBufferOrFake(LOffs.first).getBufferIdentifier(); + StringRef RB = getBufferOrFake(ROffs.first).getBufferIdentifier(); bool LIsBuiltins = LB == ""; bool RIsBuiltins = RB == ""; // Sort built-in before non-built-in. diff --git a/clang/lib/CodeGen/CGDebugInfo.cpp b/clang/lib/CodeGen/CGDebugInfo.cpp index 71ebd268f63076..8280fe718f26bd 100644 --- a/clang/lib/CodeGen/CGDebugInfo.cpp +++ b/clang/lib/CodeGen/CGDebugInfo.cpp @@ -374,9 +374,8 @@ CGDebugInfo::computeChecksum(FileID FID, SmallString<32> &Checksum) const { return None; SourceManager &SM = CGM.getContext().getSourceManager(); - bool Invalid; - const llvm::MemoryBuffer *MemBuffer = SM.getBuffer(FID, &Invalid); - if (Invalid) + Optional MemBuffer = SM.getBufferOrNone(FID); + if (!MemBuffer) return None; llvm::MD5 Hash; diff --git a/clang/lib/CodeGen/CodeGenAction.cpp b/clang/lib/CodeGen/CodeGenAction.cpp index 5ea2fc23ee1193..871cdb7dac3b9c 100644 --- a/clang/lib/CodeGen/CodeGenAction.cpp +++ b/clang/lib/CodeGen/CodeGenAction.cpp @@ -1122,11 +1122,10 @@ void CodeGenAction::ExecuteAction() { if (BA != Backend_EmitNothing && !OS) return; - bool Invalid; SourceManager &SM = CI.getSourceManager(); FileID FID = SM.getMainFileID(); - const llvm::MemoryBuffer *MainFile = SM.getBuffer(FID, &Invalid); - if (Invalid) + Optional MainFile = SM.getBufferOrNone(FID); + if (!MainFile) return; TheModule = loadModule(*MainFile); @@ -1141,8 +1140,7 @@ void CodeGenAction::ExecuteAction() { TheModule->setTargetTriple(TargetOpts.Triple); } - EmbedBitcode(TheModule.get(), CodeGenOpts, - MainFile->getMemBufferRef()); + EmbedBitcode(TheModule.get(), CodeGenOpts, *MainFile); LLVMContext &Ctx = TheModule->getContext(); Ctx.setInlineAsmDiagnosticHandler(BitcodeInlineAsmDiagHandler, diff --git a/clang/lib/Frontend/FrontendAction.cpp b/clang/lib/Frontend/FrontendAction.cpp index 92654dbe8a10a2..9a806d7c9af8d1 100644 --- a/clang/lib/Frontend/FrontendAction.cpp +++ b/clang/lib/Frontend/FrontendAction.cpp @@ -450,7 +450,7 @@ static bool loadModuleMapForModuleBuild(CompilerInstance &CI, bool IsSystem, PresumedModuleMapFile)) return true; - if (SrcMgr.getBuffer(ModuleMapID)->getBufferSize() == Offset) + if (SrcMgr.getBufferOrFake(ModuleMapID).getBufferSize() == Offset) Offset = 0; return false; diff --git a/clang/lib/Frontend/FrontendActions.cpp b/clang/lib/Frontend/FrontendActions.cpp index 77a88f696abcdd..ec5caceba20790 100644 --- a/clang/lib/Frontend/FrontendActions.cpp +++ b/clang/lib/Frontend/FrontendActions.cpp @@ -805,11 +805,9 @@ void PrintPreprocessedAction::ExecuteAction() { // concern, so if we scan for too long, we'll just assume the file should // be opened in binary mode. bool BinaryMode = true; - bool InvalidFile = false; const SourceManager& SM = CI.getSourceManager(); - const llvm::MemoryBuffer *Buffer = SM.getBuffer(SM.getMainFileID(), - &InvalidFile); - if (!InvalidFile) { + if (llvm::Optional Buffer = + SM.getBufferOrNone(SM.getMainFileID())) { const char *cur = Buffer->getBufferStart(); const char *end = Buffer->getBufferEnd(); const char *next = (cur != end) ? cur + 1 : end; @@ -937,12 +935,12 @@ void DumpCompilerOptionsAction::ExecuteAction() { void PrintDependencyDirectivesSourceMinimizerAction::ExecuteAction() { CompilerInstance &CI = getCompilerInstance(); SourceManager &SM = CI.getPreprocessor().getSourceManager(); - const llvm::MemoryBuffer *FromFile = SM.getBuffer(SM.getMainFileID()); + llvm::MemoryBufferRef FromFile = SM.getBufferOrFake(SM.getMainFileID()); llvm::SmallString<1024> Output; llvm::SmallVector Toks; if (minimizeSourceToDependencyDirectives( - FromFile->getBuffer(), Output, Toks, &CI.getDiagnostics(), + FromFile.getBuffer(), Output, Toks, &CI.getDiagnostics(), SM.getLocForStartOfFile(SM.getMainFileID()))) { assert(CI.getDiagnostics().hasErrorOccurred() && "no errors reported for failure"); diff --git a/clang/lib/Frontend/Rewrite/HTMLPrint.cpp b/clang/lib/Frontend/Rewrite/HTMLPrint.cpp index 982e56cebbca06..1388c2e1faab2b 100644 --- a/clang/lib/Frontend/Rewrite/HTMLPrint.cpp +++ b/clang/lib/Frontend/Rewrite/HTMLPrint.cpp @@ -70,7 +70,7 @@ void HTMLPrinter::HandleTranslationUnit(ASTContext &Ctx) { if (Entry) Name = Entry->getName(); else - Name = R.getSourceMgr().getBuffer(FID)->getBufferIdentifier(); + Name = R.getSourceMgr().getBufferOrFake(FID).getBufferIdentifier(); html::AddLineNumbers(R, FID); html::AddHeaderFooterInternalBuiltinCSS(R, FID, Name); diff --git a/clang/lib/Frontend/Rewrite/RewriteModernObjC.cpp b/clang/lib/Frontend/Rewrite/RewriteModernObjC.cpp index c0c81221b23444..9d5366bb161e54 100644 --- a/clang/lib/Frontend/Rewrite/RewriteModernObjC.cpp +++ b/clang/lib/Frontend/Rewrite/RewriteModernObjC.cpp @@ -702,9 +702,9 @@ void RewriteModernObjC::InitializeCommon(ASTContext &context) { // Get the ID and start/end of the main file. MainFileID = SM->getMainFileID(); - const llvm::MemoryBuffer *MainBuf = SM->getBuffer(MainFileID); - MainFileStart = MainBuf->getBufferStart(); - MainFileEnd = MainBuf->getBufferEnd(); + llvm::MemoryBufferRef MainBuf = SM->getBufferOrFake(MainFileID); + MainFileStart = MainBuf.getBufferStart(); + MainFileEnd = MainBuf.getBufferEnd(); Rewrite.setSourceMgr(Context->getSourceManager(), Context->getLangOpts()); } diff --git a/clang/lib/Frontend/Rewrite/RewriteObjC.cpp b/clang/lib/Frontend/Rewrite/RewriteObjC.cpp index 990509a84b06cf..3caf9a6720626a 100644 --- a/clang/lib/Frontend/Rewrite/RewriteObjC.cpp +++ b/clang/lib/Frontend/Rewrite/RewriteObjC.cpp @@ -631,9 +631,9 @@ void RewriteObjC::InitializeCommon(ASTContext &context) { // Get the ID and start/end of the main file. MainFileID = SM->getMainFileID(); - const llvm::MemoryBuffer *MainBuf = SM->getBuffer(MainFileID); - MainFileStart = MainBuf->getBufferStart(); - MainFileEnd = MainBuf->getBufferEnd(); + llvm::MemoryBufferRef MainBuf = SM->getBufferOrFake(MainFileID); + MainFileStart = MainBuf.getBufferStart(); + MainFileEnd = MainBuf.getBufferEnd(); Rewrite.setSourceMgr(Context->getSourceManager(), Context->getLangOpts()); } diff --git a/clang/lib/Sema/SemaExprCXX.cpp b/clang/lib/Sema/SemaExprCXX.cpp index d39820fb483d0b..8d5dccc197262e 100644 --- a/clang/lib/Sema/SemaExprCXX.cpp +++ b/clang/lib/Sema/SemaExprCXX.cpp @@ -7249,8 +7249,8 @@ ExprResult Sema::ActOnStartCXXMemberReference(Scope *S, Expr *Base, return Base; } -static bool CheckArrow(Sema& S, QualType& ObjectType, Expr *&Base, - tok::TokenKind& OpKind, SourceLocation OpLoc) { +static bool CheckArrow(Sema &S, QualType &ObjectType, Expr *&Base, + tok::TokenKind &OpKind, SourceLocation OpLoc) { if (Base->hasPlaceholderType()) { ExprResult result = S.CheckPlaceholderExpr(Base); if (result.isInvalid()) return true; @@ -7265,6 +7265,18 @@ static bool CheckArrow(Sema& S, QualType& ObjectType, Expr *&Base, // Note that this is rather different from the normal handling for the // arrow operator. if (OpKind == tok::arrow) { + // The operator requires a prvalue, so perform lvalue conversions. + // Only do this if we might plausibly end with a pointer, as otherwise + // this was likely to be intended to be a '.'. + if (ObjectType->isPointerType() || ObjectType->isArrayType() || + ObjectType->isFunctionType()) { + ExprResult BaseResult = S.DefaultFunctionArrayLvalueConversion(Base); + if (BaseResult.isInvalid()) + return true; + Base = BaseResult.get(); + ObjectType = Base->getType(); + } + if (const PointerType *Ptr = ObjectType->getAs()) { ObjectType = Ptr->getPointeeType(); } else if (!Base->isTypeDependent()) { diff --git a/clang/lib/StaticAnalyzer/Checkers/LocalizationChecker.cpp b/clang/lib/StaticAnalyzer/Checkers/LocalizationChecker.cpp index 252377f24bd7b3..28d3e058fee25b 100644 --- a/clang/lib/StaticAnalyzer/Checkers/LocalizationChecker.cpp +++ b/clang/lib/StaticAnalyzer/Checkers/LocalizationChecker.cpp @@ -1141,10 +1141,9 @@ void EmptyLocalizationContextChecker::MethodCrawler::VisitObjCMessageExpr( SE = Mgr.getSourceManager().getSLocEntry(SLInfo.first); } - bool Invalid = false; - const llvm::MemoryBuffer *BF = - Mgr.getSourceManager().getBuffer(SLInfo.first, SL, &Invalid); - if (Invalid) + llvm::Optional BF = + Mgr.getSourceManager().getBufferOrNone(SLInfo.first, SL); + if (!BF) return; Lexer TheLexer(SL, LangOptions(), BF->getBufferStart(), diff --git a/clang/lib/StaticAnalyzer/Core/BugReporter.cpp b/clang/lib/StaticAnalyzer/Core/BugReporter.cpp index 72be4e81c83d67..ebad1d1b67b429 100644 --- a/clang/lib/StaticAnalyzer/Core/BugReporter.cpp +++ b/clang/lib/StaticAnalyzer/Core/BugReporter.cpp @@ -1570,9 +1570,8 @@ static Optional getLengthOnSingleLine(const SourceManager &SM, if (FID != SM.getFileID(ExpansionRange.getEnd())) return None; - bool Invalid; - const llvm::MemoryBuffer *Buffer = SM.getBuffer(FID, &Invalid); - if (Invalid) + Optional Buffer = SM.getBufferOrNone(FID); + if (!Buffer) return None; unsigned BeginOffset = SM.getFileOffset(ExpansionRange.getBegin()); diff --git a/clang/lib/StaticAnalyzer/Core/HTMLDiagnostics.cpp b/clang/lib/StaticAnalyzer/Core/HTMLDiagnostics.cpp index bc7c41d039c4d0..68f1007569cd5f 100644 --- a/clang/lib/StaticAnalyzer/Core/HTMLDiagnostics.cpp +++ b/clang/lib/StaticAnalyzer/Core/HTMLDiagnostics.cpp @@ -786,8 +786,8 @@ void HTMLDiagnostics::HandlePiece(Rewriter &R, FileID BugFileID, if (LPosInfo.first != BugFileID) return; - const llvm::MemoryBuffer *Buf = SM.getBuffer(LPosInfo.first); - const char* FileStart = Buf->getBufferStart(); + llvm::MemoryBufferRef Buf = SM.getBufferOrFake(LPosInfo.first); + const char *FileStart = Buf.getBufferStart(); // Compute the column number. Rewind from the current position to the start // of the line. @@ -797,7 +797,7 @@ void HTMLDiagnostics::HandlePiece(Rewriter &R, FileID BugFileID, // Compute LineEnd. const char *LineEnd = TokInstantiationPtr; - const char* FileEnd = Buf->getBufferEnd(); + const char *FileEnd = Buf.getBufferEnd(); while (*LineEnd != '\n' && LineEnd != FileEnd) ++LineEnd; diff --git a/clang/lib/StaticAnalyzer/Core/IssueHash.cpp b/clang/lib/StaticAnalyzer/Core/IssueHash.cpp index e7497f3fbdaa85..2bea1593110ae2 100644 --- a/clang/lib/StaticAnalyzer/Core/IssueHash.cpp +++ b/clang/lib/StaticAnalyzer/Core/IssueHash.cpp @@ -120,7 +120,8 @@ static std::string GetEnclosingDeclContextSignature(const Decl *D) { return ""; } -static StringRef GetNthLineOfFile(const llvm::MemoryBuffer *Buffer, int Line) { +static StringRef GetNthLineOfFile(llvm::Optional Buffer, + int Line) { if (!Buffer) return ""; @@ -135,7 +136,7 @@ static std::string NormalizeLine(const SourceManager &SM, FullSourceLoc &L, const LangOptions &LangOpts) { static StringRef Whitespaces = " \t\n"; - StringRef Str = GetNthLineOfFile(SM.getBuffer(L.getFileID(), L), + StringRef Str = GetNthLineOfFile(SM.getBufferOrNone(L.getFileID(), L), L.getExpansionLineNumber()); StringRef::size_type col = Str.find_first_not_of(Whitespaces); if (col == StringRef::npos) @@ -144,8 +145,8 @@ static std::string NormalizeLine(const SourceManager &SM, FullSourceLoc &L, col++; SourceLocation StartOfLine = SM.translateLineCol(SM.getFileID(L), L.getExpansionLineNumber(), col); - const llvm::MemoryBuffer *Buffer = - SM.getBuffer(SM.getFileID(StartOfLine), StartOfLine); + Optional Buffer = + SM.getBufferOrNone(SM.getFileID(StartOfLine), StartOfLine); if (!Buffer) return {}; diff --git a/clang/lib/StaticAnalyzer/Core/PlistDiagnostics.cpp b/clang/lib/StaticAnalyzer/Core/PlistDiagnostics.cpp index 676d621e49782f..da80f154374014 100644 --- a/clang/lib/StaticAnalyzer/Core/PlistDiagnostics.cpp +++ b/clang/lib/StaticAnalyzer/Core/PlistDiagnostics.cpp @@ -887,12 +887,12 @@ class TokenStream { FileID File; unsigned Offset; std::tie(File, Offset) = SM.getDecomposedLoc(ExpanLoc); - const llvm::MemoryBuffer *MB = SM.getBuffer(File); - const char *MacroNameTokenPos = MB->getBufferStart() + Offset; + llvm::MemoryBufferRef MB = SM.getBufferOrFake(File); + const char *MacroNameTokenPos = MB.getBufferStart() + Offset; RawLexer = std::make_unique(SM.getLocForStartOfFile(File), LangOpts, - MB->getBufferStart(), MacroNameTokenPos, - MB->getBufferEnd()); + MB.getBufferStart(), MacroNameTokenPos, + MB.getBufferEnd()); } void next(Token &Result) { diff --git a/clang/lib/StaticAnalyzer/Core/SarifDiagnostics.cpp b/clang/lib/StaticAnalyzer/Core/SarifDiagnostics.cpp index 8c2e8560157684..6a9a13aed8d63e 100644 --- a/clang/lib/StaticAnalyzer/Core/SarifDiagnostics.cpp +++ b/clang/lib/StaticAnalyzer/Core/SarifDiagnostics.cpp @@ -160,9 +160,8 @@ static unsigned int adjustColumnPos(const SourceManager &SM, SourceLocation Loc, assert(LocInfo.second > SM.getExpansionColumnNumber(Loc) && "position in file is before column number?"); - bool InvalidBuffer = false; - const MemoryBuffer *Buf = SM.getBuffer(LocInfo.first, &InvalidBuffer); - assert(!InvalidBuffer && "got an invalid buffer for the location's file"); + Optional Buf = SM.getBufferOrNone(LocInfo.first); + assert(Buf && "got an invalid buffer for the location's file"); assert(Buf->getBufferSize() >= (LocInfo.second + TokenLen) && "token extends past end of buffer?"); diff --git a/clang/lib/StaticAnalyzer/Frontend/AnalysisConsumer.cpp b/clang/lib/StaticAnalyzer/Frontend/AnalysisConsumer.cpp index 392049e21c6e5b..38ed5a33a3f249 100644 --- a/clang/lib/StaticAnalyzer/Frontend/AnalysisConsumer.cpp +++ b/clang/lib/StaticAnalyzer/Frontend/AnalysisConsumer.cpp @@ -476,7 +476,7 @@ void AnalysisConsumer::HandleDeclsCallGraph(const unsigned LocalTUDeclsSize) { static bool isBisonFile(ASTContext &C) { const SourceManager &SM = C.getSourceManager(); FileID FID = SM.getMainFileID(); - StringRef Buffer = SM.getBuffer(FID)->getBuffer(); + StringRef Buffer = SM.getBufferOrFake(FID).getBuffer(); if (Buffer.startswith("/* A Bison parser, made by")) return true; return false; diff --git a/clang/test/SemaCXX/constant-expression-cxx2a.cpp b/clang/test/SemaCXX/constant-expression-cxx2a.cpp index 2aea2577d972ea..4adadc9988ab5b 100644 --- a/clang/test/SemaCXX/constant-expression-cxx2a.cpp +++ b/clang/test/SemaCXX/constant-expression-cxx2a.cpp @@ -1061,9 +1061,10 @@ namespace memory_leaks { static_assert(h({new bool(true)})); // ok } -void *operator new(std::size_t, void*); +constexpr void *operator new(std::size_t, void *p) { return p; } namespace std { template constexpr T *construct(T *p) { return new (p) T; } + template constexpr void destroy(T *p) { p->~T(); } } namespace dtor_call { @@ -1428,3 +1429,11 @@ namespace PR47805 { constexpr bool g(B b) { return &b == b.p; } static_assert(g({})); } + +constexpr bool destroy_at_test() { + int n = 0; + std::destroy(&n); + std::construct(&n); + return true; +} +static_assert(destroy_at_test()); diff --git a/clang/test/SemaCXX/pseudo-destructors.cpp b/clang/test/SemaCXX/pseudo-destructors.cpp index 7a5c540794e2dd..f214f5226ee281 100644 --- a/clang/test/SemaCXX/pseudo-destructors.cpp +++ b/clang/test/SemaCXX/pseudo-destructors.cpp @@ -183,3 +183,14 @@ namespace TwoPhaseLookup { template void f6(int *p) { p->TemplateNamesNonTemplate::C::~C(); } // expected-error {{'C' does not refer to a template}} } } + +void destroy_array_element() { + int arr[5]; + using T = int; + arr->~T(); // ok, destroy arr[0]. +} + +void destroy_function() { + using T = void(); + destroy_function->~T(); // expected-error {{object expression of non-scalar type 'void ()' cannot be used in a pseudo-destructor expression}} +} diff --git a/clang/tools/clang-diff/ClangDiff.cpp b/clang/tools/clang-diff/ClangDiff.cpp index 6d1f4b9a3ec841..3046e9c5fa371d 100644 --- a/clang/tools/clang-diff/ClangDiff.cpp +++ b/clang/tools/clang-diff/ClangDiff.cpp @@ -284,7 +284,7 @@ static unsigned printHtmlForNode(raw_ostream &OS, const diff::ASTDiff &Diff, unsigned Begin, End; std::tie(Begin, End) = Tree.getSourceRangeOffsets(Node); const SourceManager &SrcMgr = Tree.getASTContext().getSourceManager(); - auto Code = SrcMgr.getBuffer(SrcMgr.getMainFileID())->getBuffer(); + auto Code = SrcMgr.getBufferOrFake(SrcMgr.getMainFileID()).getBuffer(); for (; Offset < Begin; ++Offset) printHtml(OS, Code[Offset]); OS << "= Buffer->getBufferStart() && - LocData < Buffer->getBufferEnd()); + assert(LocData >= Buffer.getBufferStart() && + LocData < Buffer.getBufferEnd()); const char *LineBegin = LocData - LocColumn; - assert(LineBegin >= Buffer->getBufferStart()); + assert(LineBegin >= Buffer.getBufferStart()); const char *LineEnd = nullptr; for (LineEnd = LineBegin; *LineEnd != '\n' && *LineEnd != '\r' && - LineEnd < Buffer->getBufferEnd(); + LineEnd < Buffer.getBufferEnd(); ++LineEnd) ; diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp index 87138cd3b4a32d..a64f90f8f74bb9 100644 --- a/clang/tools/libclang/CIndex.cpp +++ b/clang/tools/libclang/CIndex.cpp @@ -4362,9 +4362,8 @@ const char *clang_getFileContents(CXTranslationUnit TU, CXFile file, const SourceManager &SM = cxtu::getASTUnit(TU)->getSourceManager(); FileID fid = SM.translateFile(static_cast(file)); - bool Invalid = true; - const llvm::MemoryBuffer *buf = SM.getBuffer(fid, &Invalid); - if (Invalid) { + llvm::Optional buf = SM.getBufferOrNone(fid); + if (!buf) { if (size) *size = 0; return nullptr; diff --git a/clang/unittests/AST/ASTImporterTest.cpp b/clang/unittests/AST/ASTImporterTest.cpp index 967dc035d11f32..f869e492c90ac9 100644 --- a/clang/unittests/AST/ASTImporterTest.cpp +++ b/clang/unittests/AST/ASTImporterTest.cpp @@ -5695,7 +5695,8 @@ TEST_P(ImportSourceLocations, NormalFileBuffer) { // Make sure the imported buffer has the original contents. SourceManager &ToSM = ToAST->getSourceManager(); FileID ImportedID = ToSM.getFileID(ImportedLoc); - EXPECT_EQ(Source, ToSM.getBuffer(ImportedID, SourceLocation())->getBuffer()); + EXPECT_EQ(Source, + ToSM.getBufferOrFake(ImportedID, SourceLocation()).getBuffer()); } TEST_P(ImportSourceLocations, OverwrittenFileBuffer) { @@ -5729,7 +5730,7 @@ TEST_P(ImportSourceLocations, OverwrittenFileBuffer) { SourceManager &ToSM = ToAST->getSourceManager(); FileID ImportedID = ToSM.getFileID(ImportedLoc); EXPECT_EQ(Contents, - ToSM.getBuffer(ImportedID, SourceLocation())->getBuffer()); + ToSM.getBufferOrFake(ImportedID, SourceLocation()).getBuffer()); } TEST_P(ASTImporterOptionSpecificTestBase, ImportExprOfAlignmentAttr) { diff --git a/lld/COFF/Writer.cpp b/lld/COFF/Writer.cpp index c13cd34f2ddc02..534ea82dc3c225 100644 --- a/lld/COFF/Writer.cpp +++ b/lld/COFF/Writer.cpp @@ -241,6 +241,7 @@ class Writer { void addSyntheticIdata(); void fixPartialSectionChars(StringRef name, uint32_t chars); bool fixGnuImportChunks(); + void fixTlsAlignment(); PartialSection *createPartialSection(StringRef name, uint32_t outChars); PartialSection *findPartialSection(StringRef name, uint32_t outChars); @@ -267,6 +268,7 @@ class Writer { DelayLoadContents delayIdata; EdataContents edata; bool setNoSEHCharacteristic = false; + uint32_t tlsAlignment = 0; DebugDirectoryChunk *debugDirectory = nullptr; std::vector> debugRecords; @@ -633,6 +635,11 @@ void Writer::run() { writeSections(); sortExceptionTable(); + // Fix up the alignment in the TLS Directory's characteristic field, + // if a specific alignment value is needed + if (tlsAlignment) + fixTlsAlignment(); + t1.stop(); if (!config->pdbPath.empty() && config->debug) { @@ -866,6 +873,10 @@ void Writer::createSections() { StringRef name = c->getSectionName(); if (shouldStripSectionSuffix(sc, name)) name = name.split('$').first; + + if (name.startswith(".tls")) + tlsAlignment = std::max(tlsAlignment, c->getAlignment()); + PartialSection *pSec = createPartialSection(name, c->getOutputCharacteristics()); pSec->chunks.push_back(c); @@ -2038,3 +2049,33 @@ PartialSection *Writer::findPartialSection(StringRef name, uint32_t outChars) { return it->second; return nullptr; } + +void Writer::fixTlsAlignment() { + Defined *tlsSym = + dyn_cast_or_null(symtab->findUnderscore("_tls_used")); + if (!tlsSym) + return; + + OutputSection *sec = tlsSym->getChunk()->getOutputSection(); + assert(sec && tlsSym->getRVA() >= sec->getRVA() && + "no output section for _tls_used"); + + uint8_t *secBuf = buffer->getBufferStart() + sec->getFileOff(); + uint64_t tlsOffset = tlsSym->getRVA() - sec->getRVA(); + uint64_t directorySize = config->is64() + ? sizeof(object::coff_tls_directory64) + : sizeof(object::coff_tls_directory32); + + if (tlsOffset + directorySize > sec->getRawSize()) + fatal("_tls_used sym is malformed"); + + if (config->is64()) { + object::coff_tls_directory64 *tlsDir = + reinterpret_cast(&secBuf[tlsOffset]); + tlsDir->setAlignment(tlsAlignment); + } else { + object::coff_tls_directory32 *tlsDir = + reinterpret_cast(&secBuf[tlsOffset]); + tlsDir->setAlignment(tlsAlignment); + } +} diff --git a/lld/test/COFF/tls-alignment-32.ll b/lld/test/COFF/tls-alignment-32.ll index bf9df5ca725ff7..8df11e7fd1c236 100644 --- a/lld/test/COFF/tls-alignment-32.ll +++ b/lld/test/COFF/tls-alignment-32.ll @@ -13,7 +13,8 @@ ; RUN: llvm-readobj --coff-tls-directory %t.exe | FileCheck %s ; CHECK: TLSDirectory { -; CHECK: Characteristics [ (0x0) +; CHECK: Characteristics [ (0x600000) +; CHECK-NEXT: IMAGE_SCN_ALIGN_32BYTES (0x600000) target triple = "i686-pc-windows-msvc" diff --git a/lld/test/COFF/tls-alignment-64.ll b/lld/test/COFF/tls-alignment-64.ll index 056bce97aae039..04813ca0021e7f 100644 --- a/lld/test/COFF/tls-alignment-64.ll +++ b/lld/test/COFF/tls-alignment-64.ll @@ -13,7 +13,8 @@ ; RUN: llvm-readobj --coff-tls-directory %t.exe | FileCheck %s ; CHECK: TLSDirectory { -; CHECK: Characteristics [ (0x0) +; CHECK: Characteristics [ (0x700000) +; CHECK-NEXT: IMAGE_SCN_ALIGN_64BYTES (0x700000) target triple = "x86_64-pc-windows-msvc" diff --git a/lldb/tools/debugserver/source/RNBRemote.cpp b/lldb/tools/debugserver/source/RNBRemote.cpp index b66cc8f583e8ec..aef41a5dba4381 100644 --- a/lldb/tools/debugserver/source/RNBRemote.cpp +++ b/lldb/tools/debugserver/source/RNBRemote.cpp @@ -4052,8 +4052,8 @@ rnb_err_t RNBRemote::HandlePacket_v(const char *p) { if (strcmp (err_str, "unable to start the exception thread") == 0) { snprintf (err_str, sizeof (err_str) - 1, "Not allowed to attach to process. Look in the console " - "messages (Console.app), near the debugserver entries " - "when the attached failed. The subsystem that denied " + "messages (Console.app), near the debugserver entries, " + "when the attach failed. The subsystem that denied " "the attach permission will likely have logged an " "informative message about why it was denied."); err_str[sizeof (err_str) - 1] = '\0'; diff --git a/llvm/include/llvm/DebugInfo/Symbolize/Symbolize.h b/llvm/include/llvm/DebugInfo/Symbolize/Symbolize.h index 085e4bb4ccb8f5..43c89aafea9bd0 100644 --- a/llvm/include/llvm/DebugInfo/Symbolize/Symbolize.h +++ b/llvm/include/llvm/DebugInfo/Symbolize/Symbolize.h @@ -49,6 +49,7 @@ class LLVMSymbolizer { std::string FallbackDebugPath; std::string DWPName; std::vector DebugFileDirectory; + std::function RecoverableErrorHandler = WithColor::defaultErrorHandler; }; LLVMSymbolizer() = default; diff --git a/llvm/include/llvm/Object/COFF.h b/llvm/include/llvm/Object/COFF.h index 505aab8bff5b39..e7cf1b5495c66b 100644 --- a/llvm/include/llvm/Object/COFF.h +++ b/llvm/include/llvm/Object/COFF.h @@ -576,11 +576,22 @@ struct coff_tls_directory { uint32_t getAlignment() const { // Bit [20:24] contains section alignment. - uint32_t Shift = (Characteristics & 0x00F00000) >> 20; + uint32_t Shift = (Characteristics & COFF::IMAGE_SCN_ALIGN_MASK) >> 20; if (Shift > 0) return 1U << (Shift - 1); return 0; } + + void setAlignment(uint32_t Align) { + uint32_t AlignBits = 0; + if (Align) { + assert(llvm::isPowerOf2_32(Align) && "alignment is not a power of 2"); + assert(llvm::Log2_32(Align) <= 13 && "alignment requested is too large"); + AlignBits = (llvm::Log2_32(Align) + 1) << 20; + } + Characteristics = + (Characteristics & ~COFF::IMAGE_SCN_ALIGN_MASK) | AlignBits; + } }; using coff_tls_directory32 = coff_tls_directory; diff --git a/llvm/lib/DebugInfo/Symbolize/Symbolize.cpp b/llvm/lib/DebugInfo/Symbolize/Symbolize.cpp index 383e2c0d3054c3..ed53ea9712b6da 100644 --- a/llvm/lib/DebugInfo/Symbolize/Symbolize.cpp +++ b/llvm/lib/DebugInfo/Symbolize/Symbolize.cpp @@ -68,7 +68,8 @@ LLVMSymbolizer::symbolizeCode(const ObjectFile &Obj, if (I != Modules.end()) return symbolizeCodeCommon(I->second.get(), ModuleOffset); - std::unique_ptr Context = DWARFContext::create(Obj); + std::unique_ptr Context = + DWARFContext::create(Obj, nullptr, "", Opts.RecoverableErrorHandler); Expected InfoOrErr = createModuleInfo(&Obj, std::move(Context), ModuleName); if (!InfoOrErr) @@ -572,7 +573,8 @@ LLVMSymbolizer::getOrCreateModuleInfo(const std::string &ModuleName) { } } if (!Context) - Context = DWARFContext::create(*Objects.second, nullptr, Opts.DWPName); + Context = DWARFContext::create(*Objects.second, nullptr, Opts.DWPName, + Opts.RecoverableErrorHandler); return createModuleInfo(Objects.first, std::move(Context), ModuleName); } diff --git a/llvm/lib/Target/AArch64/AArch64ISelLowering.cpp b/llvm/lib/Target/AArch64/AArch64ISelLowering.cpp index eff29d92199de3..c5af295a3cd9ce 100644 --- a/llvm/lib/Target/AArch64/AArch64ISelLowering.cpp +++ b/llvm/lib/Target/AArch64/AArch64ISelLowering.cpp @@ -12336,6 +12336,43 @@ static SDValue performSetccAddFolding(SDNode *Op, SelectionDAG &DAG) { return DAG.getNode(AArch64ISD::CSEL, dl, VT, RHS, LHS, CCVal, Cmp); } +// ADD(UADDV a, UADDV b) --> UADDV(ADD a, b) +static SDValue performUADDVCombine(SDNode *N, SelectionDAG &DAG) { + EVT VT = N->getValueType(0); + // Only scalar integer and vector types. + if (N->getOpcode() != ISD::ADD || !VT.isScalarInteger()) + return SDValue(); + + SDValue LHS = N->getOperand(0); + SDValue RHS = N->getOperand(1); + if (LHS.getOpcode() != ISD::EXTRACT_VECTOR_ELT || + RHS.getOpcode() != ISD::EXTRACT_VECTOR_ELT || LHS.getValueType() != VT) + return SDValue(); + + auto *LHSN1 = dyn_cast(LHS->getOperand(1)); + auto *RHSN1 = dyn_cast(RHS->getOperand(1)); + if (!LHSN1 || LHSN1 != RHSN1 || !RHSN1->isNullValue()) + return SDValue(); + + SDValue Op1 = LHS->getOperand(0); + SDValue Op2 = RHS->getOperand(0); + EVT OpVT1 = Op1.getValueType(); + EVT OpVT2 = Op2.getValueType(); + if (Op1.getOpcode() != AArch64ISD::UADDV || OpVT1 != OpVT2 || + Op2.getOpcode() != AArch64ISD::UADDV || + OpVT1.getVectorElementType() != VT) + return SDValue(); + + SDValue Val1 = Op1.getOperand(0); + SDValue Val2 = Op2.getOperand(0); + EVT ValVT = Val1->getValueType(0); + SDLoc DL(N); + SDValue AddVal = DAG.getNode(ISD::ADD, DL, ValVT, Val1, Val2); + return DAG.getNode(ISD::EXTRACT_VECTOR_ELT, DL, VT, + DAG.getNode(AArch64ISD::UADDV, DL, ValVT, AddVal), + DAG.getConstant(0, DL, MVT::i64)); +} + // The basic add/sub long vector instructions have variants with "2" on the end // which act on the high-half of their inputs. They are normally matched by // patterns like: @@ -12389,6 +12426,16 @@ static SDValue performAddSubLongCombine(SDNode *N, return DAG.getNode(N->getOpcode(), SDLoc(N), VT, LHS, RHS); } +static SDValue performAddSubCombine(SDNode *N, + TargetLowering::DAGCombinerInfo &DCI, + SelectionDAG &DAG) { + // Try to change sum of two reductions. + if (SDValue Val = performUADDVCombine(N, DAG)) + return Val; + + return performAddSubLongCombine(N, DCI, DAG); +} + // Massage DAGs which we can use the high-half "long" operations on into // something isel will recognize better. E.g. // @@ -14739,7 +14786,7 @@ SDValue AArch64TargetLowering::PerformDAGCombine(SDNode *N, return performABSCombine(N, DAG, DCI, Subtarget); case ISD::ADD: case ISD::SUB: - return performAddSubLongCombine(N, DCI, DAG); + return performAddSubCombine(N, DCI, DAG); case ISD::XOR: return performXorCombine(N, DAG, DCI, Subtarget); case ISD::MUL: diff --git a/llvm/lib/Target/AMDGPU/SIMemoryLegalizer.cpp b/llvm/lib/Target/AMDGPU/SIMemoryLegalizer.cpp index c03709194917a9..5786b317ea01bf 100644 --- a/llvm/lib/Target/AMDGPU/SIMemoryLegalizer.cpp +++ b/llvm/lib/Target/AMDGPU/SIMemoryLegalizer.cpp @@ -716,7 +716,7 @@ bool SIGfx6CacheControl::enableLoadCacheBypass( /// sequentially consistent, and no other thread can access scratch /// memory. - /// Other address spaces do not hava a cache. + /// Other address spaces do not have a cache. return Changed; } @@ -770,7 +770,7 @@ bool SIGfx6CacheControl::insertAcquire(MachineBasicBlock::iterator &MI, /// sequentially consistent, and no other thread can access scratch /// memory. - /// Other address spaces do not hava a cache. + /// Other address spaces do not have a cache. if (Pos == Position::AFTER) --MI; @@ -972,7 +972,7 @@ bool SIGfx10CacheControl::enableLoadCacheBypass( /// sequentially consistent, and no other thread can access scratch /// memory. - /// Other address spaces do not hava a cache. + /// Other address spaces do not have a cache. return Changed; } @@ -1035,7 +1035,7 @@ bool SIGfx10CacheControl::insertAcquire(MachineBasicBlock::iterator &MI, /// sequentially consistent, and no other thread can access scratch /// memory. - /// Other address spaces do not hava a cache. + /// Other address spaces do not have a cache. if (Pos == Position::AFTER) --MI; diff --git a/llvm/test/CodeGen/AArch64/aarch64-addv.ll b/llvm/test/CodeGen/AArch64/aarch64-addv.ll index ed3e998be0b8bc..e67e3235dd218d 100644 --- a/llvm/test/CodeGen/AArch64/aarch64-addv.ll +++ b/llvm/test/CodeGen/AArch64/aarch64-addv.ll @@ -138,11 +138,9 @@ entry: define i32 @addv_combine_i32(<4 x i32> %a1, <4 x i32> %a2) { ; CHECK-LABEL: addv_combine_i32: ; CHECK: // %bb.0: // %entry +; CHECK-NEXT: add v0.4s, v0.4s, v1.4s ; CHECK-NEXT: addv s0, v0.4s -; CHECK-NEXT: addv s1, v1.4s -; CHECK-NEXT: fmov w8, s0 -; CHECK-NEXT: fmov w9, s1 -; CHECK-NEXT: add w0, w8, w9 +; CHECK-NEXT: fmov w0, s0 ; CHECK-NEXT: ret entry: %rdx.1 = call i32 @llvm.vector.reduce.add.v4i32(<4 x i32> %a1) @@ -154,11 +152,9 @@ entry: define i64 @addv_combine_i64(<2 x i64> %a1, <2 x i64> %a2) { ; CHECK-LABEL: addv_combine_i64: ; CHECK: // %bb.0: // %entry +; CHECK-NEXT: add v0.2d, v0.2d, v1.2d ; CHECK-NEXT: addp d0, v0.2d -; CHECK-NEXT: addp d1, v1.2d -; CHECK-NEXT: fmov x8, d0 -; CHECK-NEXT: fmov x9, d1 -; CHECK-NEXT: add x0, x8, x9 +; CHECK-NEXT: fmov x0, d0 ; CHECK-NEXT: ret entry: %rdx.1 = call i64 @llvm.vector.reduce.add.v2i64(<2 x i64> %a1) diff --git a/llvm/test/CodeGen/AMDGPU/sgpr-phys-copy.mir b/llvm/test/CodeGen/AMDGPU/sgpr-phys-copy.mir new file mode 100644 index 00000000000000..306fe586faa29a --- /dev/null +++ b/llvm/test/CodeGen/AMDGPU/sgpr-phys-copy.mir @@ -0,0 +1,379 @@ +# NOTE: Assertions have been autogenerated by utils/update_mir_test_checks.py +# RUN: llc -march=amdgcn -mcpu=gfx900 -run-pass postrapseudos %s -o - | FileCheck -check-prefix=GFX9 %s + +--- +name: sgpr32 +body: | + bb.0: + liveins: $sgpr0 + ; GFX9-LABEL: name: sgpr32 + ; GFX9: $sgpr1 = S_MOV_B32 $sgpr0 + $sgpr1 = COPY $sgpr0 +... + +--- +name: sgpr32_kill +body: | + bb.0: + liveins: $sgpr0 + ; GFX9-LABEL: name: sgpr32_kill + ; GFX9: $sgpr1 = S_MOV_B32 killed $sgpr0 + $sgpr1 = COPY killed $sgpr0 +... + +--- +name: sgpr64 +body: | + bb.0: + liveins: $sgpr0_sgpr1 + ; GFX9-LABEL: name: sgpr64 + ; GFX9: $sgpr2_sgpr3 = S_MOV_B64 $sgpr0_sgpr1 + $sgpr2_sgpr3 = COPY $sgpr0_sgpr1 +... + +--- +name: sgpr64_kill +body: | + bb.0: + liveins: $sgpr0_sgpr1 + ; GFX9-LABEL: name: sgpr64_kill + ; GFX9: $sgpr2_sgpr3 = S_MOV_B64 killed $sgpr0_sgpr1 + $sgpr2_sgpr3 = COPY killed $sgpr0_sgpr1 +... + +--- +name: sgpr96_aligned_src_dst +body: | + bb.0: + liveins: $sgpr0_sgpr1_sgpr2 + ; GFX9-LABEL: name: sgpr96_aligned_src_dst + ; GFX9: $sgpr8 = S_MOV_B32 $sgpr2, implicit-def $sgpr6_sgpr7_sgpr8, implicit $sgpr0_sgpr1_sgpr2 + ; GFX9: $sgpr7 = S_MOV_B32 $sgpr1, implicit $sgpr0_sgpr1_sgpr2 + ; GFX9: $sgpr6 = S_MOV_B32 $sgpr0, implicit $sgpr0_sgpr1_sgpr2 + $sgpr6_sgpr7_sgpr8 = COPY $sgpr0_sgpr1_sgpr2 +... + +--- +name: sgpr96_aligned_src +body: | + bb.0: + liveins: $sgpr0_sgpr1_sgpr2 + ; GFX9-LABEL: name: sgpr96_aligned_src + ; GFX9: $sgpr5 = S_MOV_B32 $sgpr2, implicit-def $sgpr3_sgpr4_sgpr5, implicit $sgpr0_sgpr1_sgpr2 + ; GFX9: $sgpr4 = S_MOV_B32 $sgpr1, implicit $sgpr0_sgpr1_sgpr2 + ; GFX9: $sgpr3 = S_MOV_B32 $sgpr0, implicit $sgpr0_sgpr1_sgpr2 + $sgpr3_sgpr4_sgpr5 = COPY $sgpr0_sgpr1_sgpr2 +... + +--- +name: sgpr96_aligned_dst +body: | + bb.0: + liveins: $sgpr3_sgpr4_sgpr5 + ; GFX9-LABEL: name: sgpr96_aligned_dst + ; GFX9: $sgpr0 = S_MOV_B32 $sgpr3, implicit-def $sgpr0_sgpr1_sgpr2, implicit $sgpr3_sgpr4_sgpr5 + ; GFX9: $sgpr1 = S_MOV_B32 $sgpr4, implicit $sgpr3_sgpr4_sgpr5 + ; GFX9: $sgpr2 = S_MOV_B32 $sgpr5, implicit $sgpr3_sgpr4_sgpr5 + $sgpr0_sgpr1_sgpr2 = COPY $sgpr3_sgpr4_sgpr5 +... + +--- +name: sgpr96_unaligned_src_dst +body: | + bb.0: + liveins: $sgpr3_sgpr4_sgpr5 + ; GFX9-LABEL: name: sgpr96_unaligned_src_dst + ; GFX9: $sgpr11 = S_MOV_B32 $sgpr5, implicit-def $sgpr9_sgpr10_sgpr11, implicit $sgpr3_sgpr4_sgpr5 + ; GFX9: $sgpr10 = S_MOV_B32 $sgpr4, implicit $sgpr3_sgpr4_sgpr5 + ; GFX9: $sgpr9 = S_MOV_B32 $sgpr3, implicit $sgpr3_sgpr4_sgpr5 + $sgpr9_sgpr10_sgpr11 = COPY $sgpr3_sgpr4_sgpr5 +... + +--- +name: sgpr96_killed +body: | + bb.0: + liveins: $sgpr3_sgpr4_sgpr5 + ; GFX9-LABEL: name: sgpr96_killed + ; GFX9: $sgpr11 = S_MOV_B32 $sgpr5, implicit-def $sgpr9_sgpr10_sgpr11, implicit $sgpr3_sgpr4_sgpr5 + ; GFX9: $sgpr10 = S_MOV_B32 $sgpr4, implicit $sgpr3_sgpr4_sgpr5 + ; GFX9: $sgpr9 = S_MOV_B32 $sgpr3, implicit killed $sgpr3_sgpr4_sgpr5 + $sgpr9_sgpr10_sgpr11 = COPY killed $sgpr3_sgpr4_sgpr5 +... + +--- +name: sgpr128_forward +body: | + bb.0: + liveins: $sgpr4_sgpr5_sgpr6_sgpr7 + ; GFX9-LABEL: name: sgpr128_forward + ; GFX9: $sgpr0_sgpr1 = S_MOV_B64 $sgpr4_sgpr5, implicit-def $sgpr0_sgpr1_sgpr2_sgpr3, implicit $sgpr4_sgpr5_sgpr6_sgpr7 + ; GFX9: $sgpr2_sgpr3 = S_MOV_B64 $sgpr6_sgpr7, implicit $sgpr4_sgpr5_sgpr6_sgpr7 + $sgpr0_sgpr1_sgpr2_sgpr3 = COPY $sgpr4_sgpr5_sgpr6_sgpr7 +... + +--- +name: sgpr128_backward +body: | + bb.0: + liveins: $sgpr0_sgpr1_sgpr2_sgpr3 + ; GFX9-LABEL: name: sgpr128_backward + ; GFX9: $sgpr6_sgpr7 = S_MOV_B64 $sgpr2_sgpr3, implicit-def $sgpr4_sgpr5_sgpr6_sgpr7, implicit $sgpr0_sgpr1_sgpr2_sgpr3 + ; GFX9: $sgpr4_sgpr5 = S_MOV_B64 $sgpr0_sgpr1, implicit $sgpr0_sgpr1_sgpr2_sgpr3 + $sgpr4_sgpr5_sgpr6_sgpr7 = COPY $sgpr0_sgpr1_sgpr2_sgpr3 +... + +--- +name: sgpr128_killed +body: | + bb.0: + liveins: $sgpr4_sgpr5_sgpr6_sgpr7 + ; GFX9-LABEL: name: sgpr128_killed + ; GFX9: $sgpr0_sgpr1 = S_MOV_B64 $sgpr4_sgpr5, implicit-def $sgpr0_sgpr1_sgpr2_sgpr3, implicit $sgpr4_sgpr5_sgpr6_sgpr7 + ; GFX9: $sgpr2_sgpr3 = S_MOV_B64 $sgpr6_sgpr7, implicit killed $sgpr4_sgpr5_sgpr6_sgpr7 + $sgpr0_sgpr1_sgpr2_sgpr3 = COPY killed $sgpr4_sgpr5_sgpr6_sgpr7 +... + +--- +name: sgpr160_forward +body: | + bb.0: + liveins: $sgpr8_sgpr9_sgpr10_sgpr11_sgpr12 + ; GFX9-LABEL: name: sgpr160_forward + ; GFX9: $sgpr0 = S_MOV_B32 $sgpr8, implicit-def $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4, implicit $sgpr8_sgpr9_sgpr10_sgpr11_sgpr12 + ; GFX9: $sgpr1 = S_MOV_B32 $sgpr9, implicit $sgpr8_sgpr9_sgpr10_sgpr11_sgpr12 + ; GFX9: $sgpr2 = S_MOV_B32 $sgpr10, implicit $sgpr8_sgpr9_sgpr10_sgpr11_sgpr12 + ; GFX9: $sgpr3 = S_MOV_B32 $sgpr11, implicit $sgpr8_sgpr9_sgpr10_sgpr11_sgpr12 + ; GFX9: $sgpr4 = S_MOV_B32 $sgpr12, implicit $sgpr8_sgpr9_sgpr10_sgpr11_sgpr12 + $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4 = COPY $sgpr8_sgpr9_sgpr10_sgpr11_sgpr12 +... + +--- +name: sgpr160_backward +body: | + bb.0: + liveins: $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4 + ; GFX9-LABEL: name: sgpr160_backward + ; GFX9: $sgpr12 = S_MOV_B32 $sgpr4, implicit-def $sgpr8_sgpr9_sgpr10_sgpr11_sgpr12, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4 + ; GFX9: $sgpr11 = S_MOV_B32 $sgpr3, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4 + ; GFX9: $sgpr10 = S_MOV_B32 $sgpr2, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4 + ; GFX9: $sgpr9 = S_MOV_B32 $sgpr1, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4 + ; GFX9: $sgpr8 = S_MOV_B32 $sgpr0, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4 + $sgpr8_sgpr9_sgpr10_sgpr11_sgpr12 = COPY $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4 +... + +--- +name: sgpr160_killed +body: | + bb.0: + liveins: $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4 + ; GFX9-LABEL: name: sgpr160_killed + ; GFX9: $sgpr12 = S_MOV_B32 $sgpr4, implicit-def $sgpr8_sgpr9_sgpr10_sgpr11_sgpr12, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4 + ; GFX9: $sgpr11 = S_MOV_B32 $sgpr3, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4 + ; GFX9: $sgpr10 = S_MOV_B32 $sgpr2, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4 + ; GFX9: $sgpr9 = S_MOV_B32 $sgpr1, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4 + ; GFX9: $sgpr8 = S_MOV_B32 $sgpr0, implicit killed $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4 + $sgpr8_sgpr9_sgpr10_sgpr11_sgpr12 = COPY killed $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4 +... + + +--- +name: sgpr192_forward +body: | + bb.0: + liveins: $sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13 + ; GFX9-LABEL: name: sgpr192_forward + ; GFX9: $sgpr0_sgpr1 = S_MOV_B64 $sgpr8_sgpr9, implicit-def $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5, implicit $sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13 + ; GFX9: $sgpr2_sgpr3 = S_MOV_B64 $sgpr10_sgpr11, implicit $sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13 + ; GFX9: $sgpr4_sgpr5 = S_MOV_B64 $sgpr12_sgpr13, implicit $sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13 + $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5 = COPY $sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13 +... + +--- +name: sgpr192_backward +body: | + bb.0: + liveins: $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5 + ; GFX9-LABEL: name: sgpr192_backward + ; GFX9: $sgpr12_sgpr13 = S_MOV_B64 $sgpr4_sgpr5, implicit-def $sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5 + ; GFX9: $sgpr10_sgpr11 = S_MOV_B64 $sgpr2_sgpr3, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5 + ; GFX9: $sgpr8_sgpr9 = S_MOV_B64 $sgpr0_sgpr1, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5 + $sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13 = COPY $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5 +... + +--- +name: sgpr192_killed +body: | + bb.0: + liveins: $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5 + ; GFX9-LABEL: name: sgpr192_killed + ; GFX9: $sgpr12_sgpr13 = S_MOV_B64 $sgpr4_sgpr5, implicit-def $sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5 + ; GFX9: $sgpr10_sgpr11 = S_MOV_B64 $sgpr2_sgpr3, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5 + ; GFX9: $sgpr8_sgpr9 = S_MOV_B64 $sgpr0_sgpr1, implicit killed $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5 + $sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13 = COPY killed $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5 +... + +--- +name: sgpr256_forward +body: | + bb.0: + liveins: $sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15 + ; GFX9-LABEL: name: sgpr256_forward + ; GFX9: $sgpr0_sgpr1 = S_MOV_B64 $sgpr8_sgpr9, implicit-def $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7, implicit $sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15 + ; GFX9: $sgpr2_sgpr3 = S_MOV_B64 $sgpr10_sgpr11, implicit $sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15 + ; GFX9: $sgpr4_sgpr5 = S_MOV_B64 $sgpr12_sgpr13, implicit $sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15 + ; GFX9: $sgpr6_sgpr7 = S_MOV_B64 $sgpr14_sgpr15, implicit $sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15 + $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7 = COPY $sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15 +... + +--- +name: sgpr256_backward +body: | + bb.0: + liveins: $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7 + ; GFX9-LABEL: name: sgpr256_backward + ; GFX9: $sgpr14_sgpr15 = S_MOV_B64 $sgpr6_sgpr7, implicit-def $sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7 + ; GFX9: $sgpr12_sgpr13 = S_MOV_B64 $sgpr4_sgpr5, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7 + ; GFX9: $sgpr10_sgpr11 = S_MOV_B64 $sgpr2_sgpr3, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7 + ; GFX9: $sgpr8_sgpr9 = S_MOV_B64 $sgpr0_sgpr1, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7 + $sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15 = COPY $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7 +... + +--- +name: sgpr256_killed +body: | + bb.0: + liveins: $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7 + ; GFX9-LABEL: name: sgpr256_killed + ; GFX9: $sgpr14_sgpr15 = S_MOV_B64 $sgpr6_sgpr7, implicit-def $sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7 + ; GFX9: $sgpr12_sgpr13 = S_MOV_B64 $sgpr4_sgpr5, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7 + ; GFX9: $sgpr10_sgpr11 = S_MOV_B64 $sgpr2_sgpr3, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7 + ; GFX9: $sgpr8_sgpr9 = S_MOV_B64 $sgpr0_sgpr1, implicit killed $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7 + $sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15 = COPY killed $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7 +... + +--- +name: sgpr512_forward +body: | + bb.0: + liveins: $sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + ; GFX9-LABEL: name: sgpr512_forward + ; GFX9: $sgpr0_sgpr1 = S_MOV_B64 $sgpr16_sgpr17, implicit-def $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15, implicit $sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + ; GFX9: $sgpr2_sgpr3 = S_MOV_B64 $sgpr18_sgpr19, implicit $sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + ; GFX9: $sgpr4_sgpr5 = S_MOV_B64 $sgpr20_sgpr21, implicit $sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + ; GFX9: $sgpr6_sgpr7 = S_MOV_B64 $sgpr22_sgpr23, implicit $sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + ; GFX9: $sgpr8_sgpr9 = S_MOV_B64 $sgpr24_sgpr25, implicit $sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + ; GFX9: $sgpr10_sgpr11 = S_MOV_B64 $sgpr26_sgpr27, implicit $sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + ; GFX9: $sgpr12_sgpr13 = S_MOV_B64 $sgpr28_sgpr29, implicit $sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + ; GFX9: $sgpr14_sgpr15 = S_MOV_B64 $sgpr30_sgpr31, implicit $sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15 = COPY $sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 +... + +--- +name: sgpr512_backward +body: | + bb.0: + liveins: $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15 + ; GFX9-LABEL: name: sgpr512_backward + ; GFX9: $sgpr30_sgpr31 = S_MOV_B64 $sgpr14_sgpr15, implicit-def $sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15 + ; GFX9: $sgpr28_sgpr29 = S_MOV_B64 $sgpr12_sgpr13, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15 + ; GFX9: $sgpr26_sgpr27 = S_MOV_B64 $sgpr10_sgpr11, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15 + ; GFX9: $sgpr24_sgpr25 = S_MOV_B64 $sgpr8_sgpr9, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15 + ; GFX9: $sgpr22_sgpr23 = S_MOV_B64 $sgpr6_sgpr7, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15 + ; GFX9: $sgpr20_sgpr21 = S_MOV_B64 $sgpr4_sgpr5, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15 + ; GFX9: $sgpr18_sgpr19 = S_MOV_B64 $sgpr2_sgpr3, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15 + ; GFX9: $sgpr16_sgpr17 = S_MOV_B64 $sgpr0_sgpr1, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15 + $sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 = COPY $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15 +... + +--- +name: sgpr512_killed +body: | + bb.0: + liveins: $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15 + ; GFX9-LABEL: name: sgpr512_killed + ; GFX9: $sgpr30_sgpr31 = S_MOV_B64 $sgpr14_sgpr15, implicit-def $sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15 + ; GFX9: $sgpr28_sgpr29 = S_MOV_B64 $sgpr12_sgpr13, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15 + ; GFX9: $sgpr26_sgpr27 = S_MOV_B64 $sgpr10_sgpr11, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15 + ; GFX9: $sgpr24_sgpr25 = S_MOV_B64 $sgpr8_sgpr9, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15 + ; GFX9: $sgpr22_sgpr23 = S_MOV_B64 $sgpr6_sgpr7, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15 + ; GFX9: $sgpr20_sgpr21 = S_MOV_B64 $sgpr4_sgpr5, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15 + ; GFX9: $sgpr18_sgpr19 = S_MOV_B64 $sgpr2_sgpr3, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15 + ; GFX9: $sgpr16_sgpr17 = S_MOV_B64 $sgpr0_sgpr1, implicit killed $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15 + $sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 = COPY killed $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15 +... + +--- +name: sgpr1024_forward +body: | + bb.0: + liveins: $sgpr32_sgpr33_sgpr34_sgpr35_sgpr36_sgpr37_sgpr38_sgpr39_sgpr40_sgpr41_sgpr42_sgpr43_sgpr44_sgpr45_sgpr46_sgpr47_sgpr48_sgpr49_sgpr50_sgpr51_sgpr52_sgpr53_sgpr54_sgpr55_sgpr56_sgpr57_sgpr58_sgpr59_sgpr60_sgpr61_sgpr62_sgpr63 + ; GFX9-LABEL: name: sgpr1024_forward + ; GFX9: $sgpr0_sgpr1 = S_MOV_B64 $sgpr32_sgpr33, implicit-def $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15_sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31, implicit $sgpr32_sgpr33_sgpr34_sgpr35_sgpr36_sgpr37_sgpr38_sgpr39_sgpr40_sgpr41_sgpr42_sgpr43_sgpr44_sgpr45_sgpr46_sgpr47_sgpr48_sgpr49_sgpr50_sgpr51_sgpr52_sgpr53_sgpr54_sgpr55_sgpr56_sgpr57_sgpr58_sgpr59_sgpr60_sgpr61_sgpr62_sgpr63 + ; GFX9: $sgpr2_sgpr3 = S_MOV_B64 $sgpr34_sgpr35, implicit $sgpr32_sgpr33_sgpr34_sgpr35_sgpr36_sgpr37_sgpr38_sgpr39_sgpr40_sgpr41_sgpr42_sgpr43_sgpr44_sgpr45_sgpr46_sgpr47_sgpr48_sgpr49_sgpr50_sgpr51_sgpr52_sgpr53_sgpr54_sgpr55_sgpr56_sgpr57_sgpr58_sgpr59_sgpr60_sgpr61_sgpr62_sgpr63 + ; GFX9: $sgpr4_sgpr5 = S_MOV_B64 $sgpr36_sgpr37, implicit $sgpr32_sgpr33_sgpr34_sgpr35_sgpr36_sgpr37_sgpr38_sgpr39_sgpr40_sgpr41_sgpr42_sgpr43_sgpr44_sgpr45_sgpr46_sgpr47_sgpr48_sgpr49_sgpr50_sgpr51_sgpr52_sgpr53_sgpr54_sgpr55_sgpr56_sgpr57_sgpr58_sgpr59_sgpr60_sgpr61_sgpr62_sgpr63 + ; GFX9: $sgpr6_sgpr7 = S_MOV_B64 $sgpr38_sgpr39, implicit $sgpr32_sgpr33_sgpr34_sgpr35_sgpr36_sgpr37_sgpr38_sgpr39_sgpr40_sgpr41_sgpr42_sgpr43_sgpr44_sgpr45_sgpr46_sgpr47_sgpr48_sgpr49_sgpr50_sgpr51_sgpr52_sgpr53_sgpr54_sgpr55_sgpr56_sgpr57_sgpr58_sgpr59_sgpr60_sgpr61_sgpr62_sgpr63 + ; GFX9: $sgpr8_sgpr9 = S_MOV_B64 $sgpr40_sgpr41, implicit $sgpr32_sgpr33_sgpr34_sgpr35_sgpr36_sgpr37_sgpr38_sgpr39_sgpr40_sgpr41_sgpr42_sgpr43_sgpr44_sgpr45_sgpr46_sgpr47_sgpr48_sgpr49_sgpr50_sgpr51_sgpr52_sgpr53_sgpr54_sgpr55_sgpr56_sgpr57_sgpr58_sgpr59_sgpr60_sgpr61_sgpr62_sgpr63 + ; GFX9: $sgpr10_sgpr11 = S_MOV_B64 $sgpr42_sgpr43, implicit $sgpr32_sgpr33_sgpr34_sgpr35_sgpr36_sgpr37_sgpr38_sgpr39_sgpr40_sgpr41_sgpr42_sgpr43_sgpr44_sgpr45_sgpr46_sgpr47_sgpr48_sgpr49_sgpr50_sgpr51_sgpr52_sgpr53_sgpr54_sgpr55_sgpr56_sgpr57_sgpr58_sgpr59_sgpr60_sgpr61_sgpr62_sgpr63 + ; GFX9: $sgpr12_sgpr13 = S_MOV_B64 $sgpr44_sgpr45, implicit $sgpr32_sgpr33_sgpr34_sgpr35_sgpr36_sgpr37_sgpr38_sgpr39_sgpr40_sgpr41_sgpr42_sgpr43_sgpr44_sgpr45_sgpr46_sgpr47_sgpr48_sgpr49_sgpr50_sgpr51_sgpr52_sgpr53_sgpr54_sgpr55_sgpr56_sgpr57_sgpr58_sgpr59_sgpr60_sgpr61_sgpr62_sgpr63 + ; GFX9: $sgpr14_sgpr15 = S_MOV_B64 $sgpr46_sgpr47, implicit $sgpr32_sgpr33_sgpr34_sgpr35_sgpr36_sgpr37_sgpr38_sgpr39_sgpr40_sgpr41_sgpr42_sgpr43_sgpr44_sgpr45_sgpr46_sgpr47_sgpr48_sgpr49_sgpr50_sgpr51_sgpr52_sgpr53_sgpr54_sgpr55_sgpr56_sgpr57_sgpr58_sgpr59_sgpr60_sgpr61_sgpr62_sgpr63 + ; GFX9: $sgpr16_sgpr17 = S_MOV_B64 $sgpr48_sgpr49, implicit $sgpr32_sgpr33_sgpr34_sgpr35_sgpr36_sgpr37_sgpr38_sgpr39_sgpr40_sgpr41_sgpr42_sgpr43_sgpr44_sgpr45_sgpr46_sgpr47_sgpr48_sgpr49_sgpr50_sgpr51_sgpr52_sgpr53_sgpr54_sgpr55_sgpr56_sgpr57_sgpr58_sgpr59_sgpr60_sgpr61_sgpr62_sgpr63 + ; GFX9: $sgpr18_sgpr19 = S_MOV_B64 $sgpr50_sgpr51, implicit $sgpr32_sgpr33_sgpr34_sgpr35_sgpr36_sgpr37_sgpr38_sgpr39_sgpr40_sgpr41_sgpr42_sgpr43_sgpr44_sgpr45_sgpr46_sgpr47_sgpr48_sgpr49_sgpr50_sgpr51_sgpr52_sgpr53_sgpr54_sgpr55_sgpr56_sgpr57_sgpr58_sgpr59_sgpr60_sgpr61_sgpr62_sgpr63 + ; GFX9: $sgpr20_sgpr21 = S_MOV_B64 $sgpr52_sgpr53, implicit $sgpr32_sgpr33_sgpr34_sgpr35_sgpr36_sgpr37_sgpr38_sgpr39_sgpr40_sgpr41_sgpr42_sgpr43_sgpr44_sgpr45_sgpr46_sgpr47_sgpr48_sgpr49_sgpr50_sgpr51_sgpr52_sgpr53_sgpr54_sgpr55_sgpr56_sgpr57_sgpr58_sgpr59_sgpr60_sgpr61_sgpr62_sgpr63 + ; GFX9: $sgpr22_sgpr23 = S_MOV_B64 $sgpr54_sgpr55, implicit $sgpr32_sgpr33_sgpr34_sgpr35_sgpr36_sgpr37_sgpr38_sgpr39_sgpr40_sgpr41_sgpr42_sgpr43_sgpr44_sgpr45_sgpr46_sgpr47_sgpr48_sgpr49_sgpr50_sgpr51_sgpr52_sgpr53_sgpr54_sgpr55_sgpr56_sgpr57_sgpr58_sgpr59_sgpr60_sgpr61_sgpr62_sgpr63 + ; GFX9: $sgpr24_sgpr25 = S_MOV_B64 $sgpr56_sgpr57, implicit $sgpr32_sgpr33_sgpr34_sgpr35_sgpr36_sgpr37_sgpr38_sgpr39_sgpr40_sgpr41_sgpr42_sgpr43_sgpr44_sgpr45_sgpr46_sgpr47_sgpr48_sgpr49_sgpr50_sgpr51_sgpr52_sgpr53_sgpr54_sgpr55_sgpr56_sgpr57_sgpr58_sgpr59_sgpr60_sgpr61_sgpr62_sgpr63 + ; GFX9: $sgpr26_sgpr27 = S_MOV_B64 $sgpr58_sgpr59, implicit $sgpr32_sgpr33_sgpr34_sgpr35_sgpr36_sgpr37_sgpr38_sgpr39_sgpr40_sgpr41_sgpr42_sgpr43_sgpr44_sgpr45_sgpr46_sgpr47_sgpr48_sgpr49_sgpr50_sgpr51_sgpr52_sgpr53_sgpr54_sgpr55_sgpr56_sgpr57_sgpr58_sgpr59_sgpr60_sgpr61_sgpr62_sgpr63 + ; GFX9: $sgpr28_sgpr29 = S_MOV_B64 $sgpr60_sgpr61, implicit $sgpr32_sgpr33_sgpr34_sgpr35_sgpr36_sgpr37_sgpr38_sgpr39_sgpr40_sgpr41_sgpr42_sgpr43_sgpr44_sgpr45_sgpr46_sgpr47_sgpr48_sgpr49_sgpr50_sgpr51_sgpr52_sgpr53_sgpr54_sgpr55_sgpr56_sgpr57_sgpr58_sgpr59_sgpr60_sgpr61_sgpr62_sgpr63 + ; GFX9: $sgpr30_sgpr31 = S_MOV_B64 $sgpr62_sgpr63, implicit $sgpr32_sgpr33_sgpr34_sgpr35_sgpr36_sgpr37_sgpr38_sgpr39_sgpr40_sgpr41_sgpr42_sgpr43_sgpr44_sgpr45_sgpr46_sgpr47_sgpr48_sgpr49_sgpr50_sgpr51_sgpr52_sgpr53_sgpr54_sgpr55_sgpr56_sgpr57_sgpr58_sgpr59_sgpr60_sgpr61_sgpr62_sgpr63 + $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15_sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 = COPY $sgpr32_sgpr33_sgpr34_sgpr35_sgpr36_sgpr37_sgpr38_sgpr39_sgpr40_sgpr41_sgpr42_sgpr43_sgpr44_sgpr45_sgpr46_sgpr47_sgpr48_sgpr49_sgpr50_sgpr51_sgpr52_sgpr53_sgpr54_sgpr55_sgpr56_sgpr57_sgpr58_sgpr59_sgpr60_sgpr61_sgpr62_sgpr63 +... + +--- +name: sgpr1024_backward +body: | + bb.0: + liveins: $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15_sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + ; GFX9-LABEL: name: sgpr1024_backward + ; GFX9: $sgpr62_sgpr63 = S_MOV_B64 $sgpr30_sgpr31, implicit-def $sgpr32_sgpr33_sgpr34_sgpr35_sgpr36_sgpr37_sgpr38_sgpr39_sgpr40_sgpr41_sgpr42_sgpr43_sgpr44_sgpr45_sgpr46_sgpr47_sgpr48_sgpr49_sgpr50_sgpr51_sgpr52_sgpr53_sgpr54_sgpr55_sgpr56_sgpr57_sgpr58_sgpr59_sgpr60_sgpr61_sgpr62_sgpr63, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15_sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + ; GFX9: $sgpr60_sgpr61 = S_MOV_B64 $sgpr28_sgpr29, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15_sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + ; GFX9: $sgpr58_sgpr59 = S_MOV_B64 $sgpr26_sgpr27, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15_sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + ; GFX9: $sgpr56_sgpr57 = S_MOV_B64 $sgpr24_sgpr25, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15_sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + ; GFX9: $sgpr54_sgpr55 = S_MOV_B64 $sgpr22_sgpr23, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15_sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + ; GFX9: $sgpr52_sgpr53 = S_MOV_B64 $sgpr20_sgpr21, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15_sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + ; GFX9: $sgpr50_sgpr51 = S_MOV_B64 $sgpr18_sgpr19, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15_sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + ; GFX9: $sgpr48_sgpr49 = S_MOV_B64 $sgpr16_sgpr17, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15_sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + ; GFX9: $sgpr46_sgpr47 = S_MOV_B64 $sgpr14_sgpr15, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15_sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + ; GFX9: $sgpr44_sgpr45 = S_MOV_B64 $sgpr12_sgpr13, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15_sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + ; GFX9: $sgpr42_sgpr43 = S_MOV_B64 $sgpr10_sgpr11, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15_sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + ; GFX9: $sgpr40_sgpr41 = S_MOV_B64 $sgpr8_sgpr9, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15_sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + ; GFX9: $sgpr38_sgpr39 = S_MOV_B64 $sgpr6_sgpr7, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15_sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + ; GFX9: $sgpr36_sgpr37 = S_MOV_B64 $sgpr4_sgpr5, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15_sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + ; GFX9: $sgpr34_sgpr35 = S_MOV_B64 $sgpr2_sgpr3, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15_sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + ; GFX9: $sgpr32_sgpr33 = S_MOV_B64 $sgpr0_sgpr1, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15_sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + $sgpr32_sgpr33_sgpr34_sgpr35_sgpr36_sgpr37_sgpr38_sgpr39_sgpr40_sgpr41_sgpr42_sgpr43_sgpr44_sgpr45_sgpr46_sgpr47_sgpr48_sgpr49_sgpr50_sgpr51_sgpr52_sgpr53_sgpr54_sgpr55_sgpr56_sgpr57_sgpr58_sgpr59_sgpr60_sgpr61_sgpr62_sgpr63 = COPY $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15_sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 +... + +--- +name: sgpr1024_killed +body: | + bb.0: + liveins: $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15_sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + ; GFX9-LABEL: name: sgpr1024_killed + ; GFX9: $sgpr62_sgpr63 = S_MOV_B64 $sgpr30_sgpr31, implicit-def $sgpr32_sgpr33_sgpr34_sgpr35_sgpr36_sgpr37_sgpr38_sgpr39_sgpr40_sgpr41_sgpr42_sgpr43_sgpr44_sgpr45_sgpr46_sgpr47_sgpr48_sgpr49_sgpr50_sgpr51_sgpr52_sgpr53_sgpr54_sgpr55_sgpr56_sgpr57_sgpr58_sgpr59_sgpr60_sgpr61_sgpr62_sgpr63, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15_sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + ; GFX9: $sgpr60_sgpr61 = S_MOV_B64 $sgpr28_sgpr29, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15_sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + ; GFX9: $sgpr58_sgpr59 = S_MOV_B64 $sgpr26_sgpr27, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15_sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + ; GFX9: $sgpr56_sgpr57 = S_MOV_B64 $sgpr24_sgpr25, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15_sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + ; GFX9: $sgpr54_sgpr55 = S_MOV_B64 $sgpr22_sgpr23, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15_sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + ; GFX9: $sgpr52_sgpr53 = S_MOV_B64 $sgpr20_sgpr21, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15_sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + ; GFX9: $sgpr50_sgpr51 = S_MOV_B64 $sgpr18_sgpr19, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15_sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + ; GFX9: $sgpr48_sgpr49 = S_MOV_B64 $sgpr16_sgpr17, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15_sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + ; GFX9: $sgpr46_sgpr47 = S_MOV_B64 $sgpr14_sgpr15, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15_sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + ; GFX9: $sgpr44_sgpr45 = S_MOV_B64 $sgpr12_sgpr13, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15_sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + ; GFX9: $sgpr42_sgpr43 = S_MOV_B64 $sgpr10_sgpr11, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15_sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + ; GFX9: $sgpr40_sgpr41 = S_MOV_B64 $sgpr8_sgpr9, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15_sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + ; GFX9: $sgpr38_sgpr39 = S_MOV_B64 $sgpr6_sgpr7, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15_sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + ; GFX9: $sgpr36_sgpr37 = S_MOV_B64 $sgpr4_sgpr5, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15_sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + ; GFX9: $sgpr34_sgpr35 = S_MOV_B64 $sgpr2_sgpr3, implicit $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15_sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + ; GFX9: $sgpr32_sgpr33 = S_MOV_B64 $sgpr0_sgpr1, implicit killed $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15_sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 + $sgpr32_sgpr33_sgpr34_sgpr35_sgpr36_sgpr37_sgpr38_sgpr39_sgpr40_sgpr41_sgpr42_sgpr43_sgpr44_sgpr45_sgpr46_sgpr47_sgpr48_sgpr49_sgpr50_sgpr51_sgpr52_sgpr53_sgpr54_sgpr55_sgpr56_sgpr57_sgpr58_sgpr59_sgpr60_sgpr61_sgpr62_sgpr63 = COPY killed $sgpr0_sgpr1_sgpr2_sgpr3_sgpr4_sgpr5_sgpr6_sgpr7_sgpr8_sgpr9_sgpr10_sgpr11_sgpr12_sgpr13_sgpr14_sgpr15_sgpr16_sgpr17_sgpr18_sgpr19_sgpr20_sgpr21_sgpr22_sgpr23_sgpr24_sgpr25_sgpr26_sgpr27_sgpr28_sgpr29_sgpr30_sgpr31 +... diff --git a/llvm/test/DebugInfo/symbolize-macho-universal-unknown-arch.test b/llvm/test/DebugInfo/symbolize-macho-universal-unknown-arch.test index 744a0d3e5492a0..b6b1ee77b8b7ac 100644 --- a/llvm/test/DebugInfo/symbolize-macho-universal-unknown-arch.test +++ b/llvm/test/DebugInfo/symbolize-macho-universal-unknown-arch.test @@ -1,5 +1,5 @@ -RUN: llvm-symbolizer --obj=%p/Inputs/macho-universal 0x1f84 | FileCheck %s +RUN: not llvm-symbolizer --obj=%p/Inputs/macho-universal 0x1f84 2>&1 | FileCheck --implicit-check-not=main %s + +CHECK: LLVMSymbolizer: error reading file: Unknown architecture named: -CHECK-NOT: main CHECK: ?? -CHECK-NOT: main diff --git a/llvm/test/DebugInfo/symbolize-missing-file.test b/llvm/test/DebugInfo/symbolize-missing-file.test index 2762af4ff5dad5..bf148517f3c485 100644 --- a/llvm/test/DebugInfo/symbolize-missing-file.test +++ b/llvm/test/DebugInfo/symbolize-missing-file.test @@ -1,3 +1,3 @@ -RUN: llvm-symbolizer --obj=unexisting-file 0x1234 2>&1 | FileCheck %s +RUN: not llvm-symbolizer --obj=unexisting-file 0x1234 2>&1 | FileCheck %s CHECK: LLVMSymbolizer: error reading file: {{[Nn]}}o such file or directory diff --git a/llvm/test/tools/dsymutil/ARM/private-extern-alias.test b/llvm/test/tools/dsymutil/ARM/private-extern-alias.test new file mode 100644 index 00000000000000..99301febb15f59 --- /dev/null +++ b/llvm/test/tools/dsymutil/ARM/private-extern-alias.test @@ -0,0 +1,29 @@ +$ cat private_extern.c +__attribute__((visibility("hidden"))) +int* foo() { + int i = 10; + volatile int* j = &i; + return j; +} + +int* bar() { + return foo(); +} + +$ cat main.c +int* bar(); +int main() { + return *bar(); +} + +$ cat alias_list +_foo _baz + +$ xcrun --sdk iphoneos clang -g private_extern.c -c -o private_extern.o -target arm64-apple-ios14.0 +$ xcrun --sdk iphoneos clang -g main.c -c -o main.o -target arm64-apple-ios14.0 +$ xcrun --sdk iphoneos clang private_extern.o main.o -target arm64-apple-ios14.0 -o private_extern.out -Xlinker -alias_list -Xlinker alias_list + +RUN: dsymutil -oso-prepend-path %p/../Inputs %p/../Inputs/private/tmp/private_extern/private_extern.out -o %t.dSYM --verbose 2>&1 | FileCheck %s +CHECK-NOT: could not find object file symbol for symbol _baz +CHECK: { sym: _foo, objAddr: 0x0000000000000000, binAddr: 0x0000000100007F58, size: 0x00000020 } +CHECK: { sym: _baz, objAddr: 0x0000000000000000, binAddr: 0x0000000100007F58, size: 0x00000000 } diff --git a/llvm/test/tools/dsymutil/Inputs/private/tmp/private_extern/main.o b/llvm/test/tools/dsymutil/Inputs/private/tmp/private_extern/main.o new file mode 100644 index 00000000000000..d9cf74ee64403d Binary files /dev/null and b/llvm/test/tools/dsymutil/Inputs/private/tmp/private_extern/main.o differ diff --git a/llvm/test/tools/dsymutil/Inputs/private/tmp/private_extern/private_extern.o b/llvm/test/tools/dsymutil/Inputs/private/tmp/private_extern/private_extern.o new file mode 100644 index 00000000000000..b7f1a6a8e8472c Binary files /dev/null and b/llvm/test/tools/dsymutil/Inputs/private/tmp/private_extern/private_extern.o differ diff --git a/llvm/test/tools/dsymutil/Inputs/private/tmp/private_extern/private_extern.out b/llvm/test/tools/dsymutil/Inputs/private/tmp/private_extern/private_extern.out new file mode 100755 index 00000000000000..fff5df162cf752 Binary files /dev/null and b/llvm/test/tools/dsymutil/Inputs/private/tmp/private_extern/private_extern.out differ diff --git a/llvm/test/tools/llvm-dwarfdump/cmdline.test b/llvm/test/tools/llvm-dwarfdump/cmdline.test index 5c7b998338c72d..11f6b15d450854 100644 --- a/llvm/test/tools/llvm-dwarfdump/cmdline.test +++ b/llvm/test/tools/llvm-dwarfdump/cmdline.test @@ -31,7 +31,7 @@ HELP: @FILE RUN: llvm-dwarfdump --version 2>&1 | FileCheck --check-prefix=VERSION %s VERSION: {{ version }} -RUN: llvm-dwarfdump -diff -verbose 2>&1 | FileCheck --check-prefix=INCOMPATIBLE %s +RUN: not llvm-dwarfdump -diff -verbose 2>&1 | FileCheck --check-prefix=INCOMPATIBLE %s INCOMPATIBLE: error: incompatible arguments: specifying both -diff and -verbose is currently not supported RUN: not llvm-dwarfdump --debug-names=0x0 2>&1 | FileCheck --check-prefix=FLAG %s diff --git a/llvm/test/tools/llvm-symbolizer/input-base.test b/llvm/test/tools/llvm-symbolizer/input-base.test index 66244a7203c088..f4b641ea626f21 100644 --- a/llvm/test/tools/llvm-symbolizer/input-base.test +++ b/llvm/test/tools/llvm-symbolizer/input-base.test @@ -1,33 +1,34 @@ -# llvm-symbolizer infers the number base from the form of the address. -RUN: llvm-symbolizer -e /dev/null -a 0x1234 | FileCheck %s -RUN: llvm-symbolizer -e /dev/null -a 0X1234 | FileCheck %s -RUN: llvm-symbolizer -e /dev/null -a 4660 | FileCheck %s -RUN: llvm-symbolizer -e /dev/null -a 011064 | FileCheck %s -RUN: llvm-symbolizer -e /dev/null -a 0b1001000110100 | FileCheck %s -RUN: llvm-symbolizer -e /dev/null -a 0B1001000110100 | FileCheck %s -RUN: llvm-symbolizer -e /dev/null -a 0o11064 | FileCheck %s +# RUN: llvm-mc %s -o %t -filetype=obj +## llvm-symbolizer infers the number base from the form of the address. +# RUN: llvm-symbolizer -e %t -a 0x1234 | FileCheck %s +# RUN: llvm-symbolizer -e %t -a 0X1234 | FileCheck %s +# RUN: llvm-symbolizer -e %t -a 4660 | FileCheck %s +# RUN: llvm-symbolizer -e %t -a 011064 | FileCheck %s +# RUN: llvm-symbolizer -e %t -a 0b1001000110100 | FileCheck %s +# RUN: llvm-symbolizer -e %t -a 0B1001000110100 | FileCheck %s +# RUN: llvm-symbolizer -e %t -a 0o11064 | FileCheck %s -# llvm-symbolizer / StringRef::getAsInteger only accepts the 0o prefix in lowercase. -RUN: llvm-symbolizer -e /dev/null -a 0O1234 | FileCheck %s --check-prefix=INVALID-NOT-OCTAL-UPPER +## llvm-symbolizer / StringRef::getAsInteger only accepts the 0o prefix in lowercase. +# RUN: llvm-symbolizer -e %t -a 0O1234 | FileCheck %s --check-prefix=INVALID-NOT-OCTAL-UPPER -# llvm-addr2line always requires hexadecimal, but accepts an optional 0x prefix. -RUN: llvm-addr2line -e /dev/null -a 0x1234 | FileCheck %s -RUN: llvm-addr2line -e /dev/null -a 0X1234 | FileCheck %s -RUN: llvm-addr2line -e /dev/null -a 1234 | FileCheck %s -RUN: llvm-addr2line -e /dev/null -a 01234 | FileCheck %s -RUN: llvm-addr2line -e /dev/null -a 0b1010 | FileCheck %s --check-prefix=HEXADECIMAL-NOT-BINARY -RUN: llvm-addr2line -e /dev/null -a 0B1010 | FileCheck %s --check-prefix=HEXADECIMAL-NOT-BINARY -RUN: llvm-addr2line -e /dev/null -a 0o1234 | FileCheck %s --check-prefix=INVALID-NOT-OCTAL-LOWER -RUN: llvm-addr2line -e /dev/null -a 0O1234 | FileCheck %s --check-prefix=INVALID-NOT-OCTAL-UPPER +## llvm-addr2line always requires hexadecimal, but accepts an optional 0x prefix. +# RUN: llvm-addr2line -e %t -a 0x1234 | FileCheck %s +# RUN: llvm-addr2line -e %t -a 0X1234 | FileCheck %s +# RUN: llvm-addr2line -e %t -a 1234 | FileCheck %s +# RUN: llvm-addr2line -e %t -a 01234 | FileCheck %s +# RUN: llvm-addr2line -e %t -a 0b1010 | FileCheck %s --check-prefix=HEXADECIMAL-NOT-BINARY +# RUN: llvm-addr2line -e %t -a 0B1010 | FileCheck %s --check-prefix=HEXADECIMAL-NOT-BINARY +# RUN: llvm-addr2line -e %t -a 0o1234 | FileCheck %s --check-prefix=INVALID-NOT-OCTAL-LOWER +# RUN: llvm-addr2line -e %t -a 0O1234 | FileCheck %s --check-prefix=INVALID-NOT-OCTAL-UPPER -CHECK: 0x1234 -CHECK-NEXT: ?? +# CHECK: 0x1234 +# CHECK-NEXT: ?? -HEXADECIMAL-NOT-BINARY: 0xb1010 -HEXADECIMAL-NOT-BINARY: ?? +# HEXADECIMAL-NOT-BINARY: 0xb1010 +# HEXADECIMAL-NOT-BINARY: ?? -INVALID-NOT-OCTAL-LOWER: 0o1234 -INVALID-NOT-OCTAL-LOWER-NOT: ?? +# INVALID-NOT-OCTAL-LOWER: 0o1234 +# INVALID-NOT-OCTAL-LOWER-NOT: ?? -INVALID-NOT-OCTAL-UPPER: 0O1234 -INVALID-NOT-OCTAL-UPPER-NOT: ?? +# INVALID-NOT-OCTAL-UPPER: 0O1234 +# INVALID-NOT-OCTAL-UPPER-NOT: ?? diff --git a/llvm/test/tools/llvm-symbolizer/invalid-dwarf.s b/llvm/test/tools/llvm-symbolizer/invalid-dwarf.s new file mode 100644 index 00000000000000..d9039acb245746 --- /dev/null +++ b/llvm/test/tools/llvm-symbolizer/invalid-dwarf.s @@ -0,0 +1,204 @@ +# Source: +# void f1() { } +# void f2() { } +# +# Build as: clang -ffunction-sections -gdwarf-5 -c test.c +# Hand modify the rnglist to include an invalid RLE encoding (42) +# +# RUN: llvm-mc -dwarf-version=5 %s -filetype=obj -o %t +# RUN: not llvm-symbolizer -obj=%t 0x0 2>&1 | FileCheck %s +# +# CHECK: error: decoding address ranges: unknown rnglists encoding 0x2a at offset 0x10 + + .text + .file "test.c" + .section .text.f1,"ax",@progbits + .globl f1 # -- Begin function f1 + .p2align 4, 0x90 + .type f1,@function +f1: # @f1 +.Lfunc_begin0: + .file 0 "/usr/local/google/home/blaikie/dev/scratch" "test.c" md5 0x39c2464ceaf7fd68a00d44ca40e99028 + .loc 0 1 0 # test.c:1:0 + .cfi_startproc +# %bb.0: # %entry + pushq %rbp + .cfi_def_cfa_offset 16 + .cfi_offset %rbp, -16 + movq %rsp, %rbp + .cfi_def_cfa_register %rbp +.Ltmp0: + .loc 0 1 13 prologue_end # test.c:1:13 + popq %rbp + .cfi_def_cfa %rsp, 8 + retq +.Ltmp1: +.Lfunc_end0: + .size f1, .Lfunc_end0-f1 + .cfi_endproc + # -- End function + .section .text.f2,"ax",@progbits + .globl f2 # -- Begin function f2 + .p2align 4, 0x90 + .type f2,@function +f2: # @f2 +.Lfunc_begin1: + .loc 0 2 0 # test.c:2:0 + .cfi_startproc +# %bb.0: # %entry + pushq %rbp + .cfi_def_cfa_offset 16 + .cfi_offset %rbp, -16 + movq %rsp, %rbp + .cfi_def_cfa_register %rbp +.Ltmp2: + .loc 0 2 13 prologue_end # test.c:2:13 + popq %rbp + .cfi_def_cfa %rsp, 8 + retq +.Ltmp3: +.Lfunc_end1: + .size f2, .Lfunc_end1-f2 + .cfi_endproc + # -- End function + .section .debug_abbrev,"",@progbits + .byte 1 # Abbreviation Code + .byte 17 # DW_TAG_compile_unit + .byte 1 # DW_CHILDREN_yes + .byte 37 # DW_AT_producer + .byte 37 # DW_FORM_strx1 + .byte 19 # DW_AT_language + .byte 5 # DW_FORM_data2 + .byte 3 # DW_AT_name + .byte 37 # DW_FORM_strx1 + .byte 114 # DW_AT_str_offsets_base + .byte 23 # DW_FORM_sec_offset + .byte 16 # DW_AT_stmt_list + .byte 23 # DW_FORM_sec_offset + .byte 27 # DW_AT_comp_dir + .byte 37 # DW_FORM_strx1 + .byte 17 # DW_AT_low_pc + .byte 1 # DW_FORM_addr + .byte 85 # DW_AT_ranges + .byte 35 # DW_FORM_rnglistx + .byte 115 # DW_AT_addr_base + .byte 23 # DW_FORM_sec_offset + .byte 116 # DW_AT_rnglists_base + .byte 23 # DW_FORM_sec_offset + .byte 0 # EOM(1) + .byte 0 # EOM(2) + .byte 2 # Abbreviation Code + .byte 46 # DW_TAG_subprogram + .byte 0 # DW_CHILDREN_no + .byte 17 # DW_AT_low_pc + .byte 27 # DW_FORM_addrx + .byte 18 # DW_AT_high_pc + .byte 6 # DW_FORM_data4 + .byte 64 # DW_AT_frame_base + .byte 24 # DW_FORM_exprloc + .byte 3 # DW_AT_name + .byte 37 # DW_FORM_strx1 + .byte 58 # DW_AT_decl_file + .byte 11 # DW_FORM_data1 + .byte 59 # DW_AT_decl_line + .byte 11 # DW_FORM_data1 + .byte 63 # DW_AT_external + .byte 25 # DW_FORM_flag_present + .byte 0 # EOM(1) + .byte 0 # EOM(2) + .byte 0 # EOM(3) + .section .debug_info,"",@progbits +.Lcu_begin0: + .long .Ldebug_info_end0-.Ldebug_info_start0 # Length of Unit +.Ldebug_info_start0: + .short 5 # DWARF version number + .byte 1 # DWARF Unit Type + .byte 8 # Address Size (in bytes) + .long .debug_abbrev # Offset Into Abbrev. Section + .byte 1 # Abbrev [1] 0xc:0x36 DW_TAG_compile_unit + .byte 0 # DW_AT_producer + .short 12 # DW_AT_language + .byte 1 # DW_AT_name + .long .Lstr_offsets_base0 # DW_AT_str_offsets_base + .long .Lline_table_start0 # DW_AT_stmt_list + .byte 2 # DW_AT_comp_dir + .quad 0 # DW_AT_low_pc + .byte 0 # DW_AT_ranges + .long .Laddr_table_base0 # DW_AT_addr_base + .long .Lrnglists_table_base0 # DW_AT_rnglists_base + .byte 2 # Abbrev [2] 0x2b:0xb DW_TAG_subprogram + .byte 0 # DW_AT_low_pc + .long .Lfunc_end0-.Lfunc_begin0 # DW_AT_high_pc + .byte 1 # DW_AT_frame_base + .byte 86 + .byte 3 # DW_AT_name + .byte 0 # DW_AT_decl_file + .byte 1 # DW_AT_decl_line + # DW_AT_external + .byte 2 # Abbrev [2] 0x36:0xb DW_TAG_subprogram + .byte 1 # DW_AT_low_pc + .long .Lfunc_end1-.Lfunc_begin1 # DW_AT_high_pc + .byte 1 # DW_AT_frame_base + .byte 86 + .byte 4 # DW_AT_name + .byte 0 # DW_AT_decl_file + .byte 2 # DW_AT_decl_line + # DW_AT_external + .byte 0 # End Of Children Mark +.Ldebug_info_end0: + .section .debug_rnglists,"",@progbits + .long .Ldebug_list_header_end0-.Ldebug_list_header_start0 # Length +.Ldebug_list_header_start0: + .short 5 # Version + .byte 8 # Address size + .byte 0 # Segment selector size + .long 1 # Offset entry count +.Lrnglists_table_base0: + .long .Ldebug_ranges0-.Lrnglists_table_base0 +.Ldebug_ranges0: + .byte 42 # DW_RLE_startx_length + .byte 0 # start index + .uleb128 .Lfunc_end0-.Lfunc_begin0 # length + .byte 3 # DW_RLE_startx_length + .byte 1 # start index + .uleb128 .Lfunc_end1-.Lfunc_begin1 # length + .byte 0 # DW_RLE_end_of_list +.Ldebug_list_header_end0: + .section .debug_str_offsets,"",@progbits + .long 24 # Length of String Offsets Set + .short 5 + .short 0 +.Lstr_offsets_base0: + .section .debug_str,"MS",@progbits,1 +.Linfo_string0: + .asciz "clang version 12.0.0 (git@github.com:llvm/llvm-project.git 9a33f027ac7d73e14ae287e78ab554142d1cbc8f)" # string offset=0 +.Linfo_string1: + .asciz "test.c" # string offset=101 +.Linfo_string2: + .asciz "/usr/local/google/home/blaikie/dev/scratch" # string offset=108 +.Linfo_string3: + .asciz "f1" # string offset=151 +.Linfo_string4: + .asciz "f2" # string offset=154 + .section .debug_str_offsets,"",@progbits + .long .Linfo_string0 + .long .Linfo_string1 + .long .Linfo_string2 + .long .Linfo_string3 + .long .Linfo_string4 + .section .debug_addr,"",@progbits + .long .Ldebug_addr_end0-.Ldebug_addr_start0 # Length of contribution +.Ldebug_addr_start0: + .short 5 # DWARF version number + .byte 8 # Address size + .byte 0 # Segment selector size +.Laddr_table_base0: + .quad .Lfunc_begin0 + .quad .Lfunc_begin1 +.Ldebug_addr_end0: + .ident "clang version 12.0.0 (git@github.com:llvm/llvm-project.git 9a33f027ac7d73e14ae287e78ab554142d1cbc8f)" + .section ".note.GNU-stack","",@progbits + .addrsig + .section .debug_line,"",@progbits +.Lline_table_start0: + diff --git a/llvm/test/tools/llvm-symbolizer/sym.test b/llvm/test/tools/llvm-symbolizer/sym.test index 73097b2c4b9b6a..612485b8b1fe64 100644 --- a/llvm/test/tools/llvm-symbolizer/sym.test +++ b/llvm/test/tools/llvm-symbolizer/sym.test @@ -31,7 +31,8 @@ RUN: llvm-symbolizer -i -print-address -p -obj=%p/Inputs/addr.exe < %p/Inputs/ad RUN: llvm-symbolizer --inlining=true --print-address -p --obj=%p/Inputs/addr.exe < %p/Inputs/addr.inp | FileCheck -check-prefix="PRETTY" %s RUN: echo "0x1" > %t.input -RUN: llvm-symbolizer -obj=%p/Inputs/zero < %t.input | FileCheck -check-prefix="ZERO" %s +## FIXME: Looks like this was meant to have an input file to test, but it isn't present - so this test probably isn't testing what it's intended to test +RUN: not llvm-symbolizer -obj=%p/Inputs/zero < %t.input | FileCheck -check-prefix="ZERO" %s RUN: llvm-addr2line -obj=%p/Inputs/addr.exe < %p/Inputs/addr.inp | FileCheck -check-prefix=A2L %s RUN: llvm-addr2line -a -obj=%p/Inputs/addr.exe < %p/Inputs/addr.inp | FileCheck -check-prefixes=A2L,A2L_A %s diff --git a/llvm/tools/dsymutil/MachODebugMapParser.cpp b/llvm/tools/dsymutil/MachODebugMapParser.cpp index 617b708561c42d..3323a0514e324a 100644 --- a/llvm/tools/dsymutil/MachODebugMapParser.cpp +++ b/llvm/tools/dsymutil/MachODebugMapParser.cpp @@ -562,7 +562,9 @@ void MachODebugMapParser::loadMainBinarySymbols( continue; } Section = *SectionOrErr; - if (Section == MainBinary.section_end() || Section->isText()) + if ((Section == MainBinary.section_end() || Section->isText()) && + !(SymType & + MachO::N_PEXT)) // Alias to non-external (was a private external) continue; uint64_t Addr = cantFail(Sym.getValue()); Expected NameOrErr = Sym.getName(); diff --git a/llvm/tools/llvm-dwarfdump/llvm-dwarfdump.cpp b/llvm/tools/llvm-dwarfdump/llvm-dwarfdump.cpp index d8fa4f9953dcad..abaf3140bcfb9c 100644 --- a/llvm/tools/llvm-dwarfdump/llvm-dwarfdump.cpp +++ b/llvm/tools/llvm-dwarfdump/llvm-dwarfdump.cpp @@ -624,7 +624,7 @@ int main(int argc, char **argv) { if (Diff && Verbose) { WithColor::error() << "incompatible arguments: specifying both -diff and " "-verbose is currently not supported"; - return 0; + return 1; } std::error_code EC; diff --git a/llvm/tools/llvm-symbolizer/llvm-symbolizer.cpp b/llvm/tools/llvm-symbolizer/llvm-symbolizer.cpp index f57922b2c506e1..e4061408a821d5 100644 --- a/llvm/tools/llvm-symbolizer/llvm-symbolizer.cpp +++ b/llvm/tools/llvm-symbolizer/llvm-symbolizer.cpp @@ -73,12 +73,15 @@ static cl::list ClInputAddresses(cl::Positional, cl::desc("..."), cl::ZeroOrMore); +static bool HasError = false; + template static bool error(Expected &ResOrErr) { if (ResOrErr) return false; logAllUnhandledErrors(ResOrErr.takeError(), errs(), "LLVMSymbolizer: error reading file: "); + HasError = true; return true; } @@ -290,6 +293,10 @@ int main(int argc, char **argv) { Args.hasFlag(OPT_untag_addresses, OPT_no_untag_addresses, !IsAddr2Line); Opts.UseNativePDBReader = Args.hasArg(OPT_use_native_pdb_reader); Opts.UseSymbolTable = true; + Opts.RecoverableErrorHandler = [&](Error E) { + HasError = true; + WithColor::defaultErrorHandler(std::move(E)); + }; for (const opt::Arg *A : Args.filtered(OPT_dsym_hint_EQ)) { StringRef Hint(A->getValue()); @@ -336,5 +343,5 @@ int main(int argc, char **argv) { Symbolizer, Printer); } - return 0; + return HasError; } diff --git a/mlir/include/mlir/Dialect/SPIRV/TargetAndABI.td b/mlir/include/mlir/Dialect/SPIRV/TargetAndABI.td index e8b1665410e5d5..13bffabaef1798 100644 --- a/mlir/include/mlir/Dialect/SPIRV/TargetAndABI.td +++ b/mlir/include/mlir/Dialect/SPIRV/TargetAndABI.td @@ -36,6 +36,25 @@ def SPV_ExtensionArrayAttr : TypedArrayAttrBase< def SPV_CapabilityArrayAttr : TypedArrayAttrBase< SPV_CapabilityAttr, "SPIR-V capability array attribute">; +// Description of cooperative matrix operations supported on the +// target. Represents `VkCooperativeMatrixPropertiesNV`. See +// https://www.khronos.org/registry/vulkan/specs/1.2-extensions/man/html/VkCooperativeMatrixPropertiesNV.html +def SPV_CooperativeMatrixPropertiesNVAttr : + StructAttr<"CooperativeMatrixPropertiesNVAttr", SPIRV_Dialect, [ + StructFieldAttr<"m_size", I32Attr>, + StructFieldAttr<"n_size", I32Attr>, + StructFieldAttr<"k_size", I32Attr>, + StructFieldAttr<"a_type", TypeAttr>, + StructFieldAttr<"b_type", TypeAttr>, + StructFieldAttr<"c_type", TypeAttr>, + StructFieldAttr<"result_type", TypeAttr>, + StructFieldAttr<"scope", SPV_ScopeAttr> +]>; + +def SPV_CooperativeMatrixPropertiesNVArrayAttr : + TypedArrayAttrBase; + // This attribute specifies the limits for various resources on the target // architecture. // @@ -60,7 +79,13 @@ def SPV_ResourceLimitsAttr : StructAttr<"ResourceLimitsAttr", SPIRV_Dialect, [ // The default number of invocations in each subgroup. // 0x7FFFFFFF means unknown. - StructFieldAttr<"subgroup_size", DefaultValuedAttr> + StructFieldAttr<"subgroup_size", DefaultValuedAttr>, + + // The configurations of cooperative matrix operations + // supported. Default is an empty list. + StructFieldAttr< + "cooperative_matrix_properties_nv", + DefaultValuedAttr> ]>; #endif // SPIRV_TARGET_AND_ABI diff --git a/mlir/lib/Dialect/SPIRV/TargetAndABI.cpp b/mlir/lib/Dialect/SPIRV/TargetAndABI.cpp index ae076513f03100..29a89a30ad4cf2 100644 --- a/mlir/lib/Dialect/SPIRV/TargetAndABI.cpp +++ b/mlir/lib/Dialect/SPIRV/TargetAndABI.cpp @@ -140,7 +140,8 @@ spirv::getDefaultResourceLimits(MLIRContext *context) { /*max_compute_shared_memory_size=*/nullptr, /*max_compute_workgroup_invocations=*/nullptr, /*max_compute_workgroup_size=*/nullptr, - /*subgroup_size=*/nullptr, context); + /*subgroup_size=*/nullptr, + /*cooperative_matrix_properties_nv=*/nullptr, context); } StringRef spirv::getTargetEnvAttrName() { return "spv.target_env"; } diff --git a/mlir/test/Dialect/SPIRV/target-and-abi.mlir b/mlir/test/Dialect/SPIRV/target-and-abi.mlir index 6edc9172677869..5f51680e97ea26 100644 --- a/mlir/test/Dialect/SPIRV/target-and-abi.mlir +++ b/mlir/test/Dialect/SPIRV/target-and-abi.mlir @@ -171,6 +171,44 @@ func @target_env_extra_fields() attributes { // ----- +func @target_env_cooperative_matrix() attributes{ + // CHECK: spv.target_env = #spv.target_env< + // CHECK-SAME: SPV_NV_cooperative_matrix + // CHECK-SAME: cooperative_matrix_properties_nv = [ + // CHECK-SAME: {a_type = i8, b_type = i8, c_type = i32, + // CHECK-SAME: k_size = 32 : i32, m_size = 8 : i32, n_size = 8 : i32 + // CHECK-SAME: result_type = i32, scope = 3 : i32} + // CHECK-SAME: {a_type = f16, b_type = f16, c_type = f16, + // CHECK-SAME: k_size = 16 : i32, m_size = 8 : i32, n_size = 8 : i32 + // CHECK-SAME: result_type = f16, scope = 3 : i32} + spv.target_env = #spv.target_env< + #spv.vce, + { + cooperative_matrix_properties_nv = [{ + m_size = 8: i32, + n_size = 8: i32, + k_size = 32: i32, + a_type = i8, + b_type = i8, + c_type = i32, + result_type = i32, + scope = 3: i32 + }, { + m_size = 8: i32, + n_size = 8: i32, + k_size = 16: i32, + a_type = f16, + b_type = f16, + c_type = f16, + result_type = f16, + scope = 3: i32 + }] + }> +} { return } + +// ----- + //===----------------------------------------------------------------------===// // spv.vce //===----------------------------------------------------------------------===//