[llvm] Extend llvm objdump fatbin (PR #114834)
David Salinas via llvm-commits
llvm-commits at lists.llvm.org
Mon Nov 11 12:12:09 PST 2024
https://github.com/david-salinas updated https://github.com/llvm/llvm-project/pull/114834
>From b6e9520a591d451378bac3447d8c41d632f6923a Mon Sep 17 00:00:00 2001
From: David <dsalinas at amd.com>
Date: Wed, 28 Aug 2024 04:01:38 +0100
Subject: [PATCH 1/3] Extend LLVM Offloading API for binary fatbin Bundles
With the intention to provide a common API for offloading, this
extension to the existing LLVM Offloading API adds support for
Binary Fatbin Bundles; moving some support from the Clang offloading
API. The intention is to add functionality to LLVM tooling for
Binary Fatbin Bundles in subsequent commits.
Change-Id: I907fdcbcd0545162a0ce1cf17ebf7c9f3a4dbde6
---
llvm/include/llvm/Object/OffloadBinary.h | 153 ++++++++
llvm/lib/Object/OffloadBinary.cpp | 468 +++++++++++++++++++++++
2 files changed, 621 insertions(+)
diff --git a/llvm/include/llvm/Object/OffloadBinary.h b/llvm/include/llvm/Object/OffloadBinary.h
index c02aec8d956ed6..ee01cabe146517 100644
--- a/llvm/include/llvm/Object/OffloadBinary.h
+++ b/llvm/include/llvm/Object/OffloadBinary.h
@@ -21,6 +21,8 @@
#include "llvm/ADT/SmallString.h"
#include "llvm/ADT/StringRef.h"
#include "llvm/Object/Binary.h"
+#include "llvm/Object/ObjectFile.h"
+#include "llvm/Support/Compression.h"
#include "llvm/Support/Error.h"
#include "llvm/Support/MemoryBuffer.h"
#include <memory>
@@ -49,6 +51,31 @@ enum ImageKind : uint16_t {
IMG_LAST,
};
+class CompressedOffloadBundle {
+private:
+ static inline const size_t MagicSize = 4;
+ static inline const size_t VersionFieldSize = sizeof(uint16_t);
+ static inline const size_t MethodFieldSize = sizeof(uint16_t);
+ static inline const size_t FileSizeFieldSize = sizeof(uint32_t);
+ static inline const size_t UncompressedSizeFieldSize = sizeof(uint32_t);
+ static inline const size_t HashFieldSize = sizeof(uint64_t);
+ static inline const size_t V1HeaderSize =
+ MagicSize + VersionFieldSize + MethodFieldSize +
+ UncompressedSizeFieldSize + HashFieldSize;
+ static inline const size_t V2HeaderSize =
+ MagicSize + VersionFieldSize + FileSizeFieldSize + MethodFieldSize +
+ UncompressedSizeFieldSize + HashFieldSize;
+ static inline const llvm::StringRef MagicNumber = "CCOB";
+ static inline const uint16_t Version = 2;
+
+public:
+ static llvm::Expected<std::unique_ptr<llvm::MemoryBuffer>>
+ compress(llvm::compression::Params P, const llvm::MemoryBuffer &Input,
+ bool Verbose = false);
+ static llvm::Expected<std::unique_ptr<llvm::MemoryBuffer>>
+ decompress(llvm::MemoryBufferRef &Input, bool Verbose = false);
+};
+
/// A simple binary serialization of an offloading file. We use this format to
/// embed the offloading image into the host executable so it can be extracted
/// and used by the linker.
@@ -183,11 +210,137 @@ class OffloadFile : public OwningBinary<OffloadBinary> {
}
};
+class OffloadFatBinBundle {
+
+private:
+ uint64_t Size = 0u;
+ StringRef FileName;
+ int64_t NumberOfEntries;
+
+public:
+ struct BundleEntry {
+ uint64_t Offset = 0u;
+ uint64_t Size = 0u;
+ uint64_t IDLength = 0u;
+ StringRef ID;
+ BundleEntry(uint64_t O, uint64_t S, uint64_t I, StringRef T)
+ : Offset(O), Size(S), IDLength(I), ID(T) {}
+ void dump(raw_ostream &OS) {
+ OS << "Offset = " << Offset << ", Size = " << Size
+ << ", ID Length = " << IDLength << ", ID = " << ID;
+ }
+ void dumpURI(raw_ostream &OS, StringRef filePath) {
+ OS << ID.data() << "\tfile:\/\/" << filePath << "#offset=" << Offset
+ << "&size=" << Size << "\n";
+ }
+ };
+
+ uint64_t getSize() const { return Size; }
+ StringRef getFileName() const { return FileName; }
+ int64_t getNumEntries() const { return NumberOfEntries; }
+
+ std::unique_ptr<SmallVector<BundleEntry>> Entries;
+ static Expected<std::unique_ptr<OffloadFatBinBundle>>
+ create(MemoryBufferRef, uint64_t SectionOffset, StringRef fileName);
+ Error extractBundle(const ObjectFile &Source);
+
+ Error ReadEntries(StringRef Section, uint64_t SectionOffset);
+ void DumpEntries() {
+ SmallVectorImpl<BundleEntry>::iterator it = Entries->begin();
+ for (int64_t I = 0; I < Entries->size(); I++) {
+ it->dump(outs());
+ ++it;
+ }
+ }
+
+ void PrintEntriesAsURI() {
+ SmallVectorImpl<BundleEntry>::iterator it = Entries->begin();
+ for (int64_t I = 0; I < NumberOfEntries; I++) {
+ it->dumpURI(outs(), FileName);
+ ++it;
+ }
+ }
+
+ OffloadFatBinBundle(MemoryBufferRef Source, StringRef file) : FileName(file) {
+ NumberOfEntries = 0;
+ Entries = std::make_unique<SmallVector<BundleEntry>>();
+ }
+};
+
+enum uri_type_t { FILE_URI, MEMORY_URI };
+
+struct OffloadBundleURI {
+ int64_t Offset = 0;
+ int64_t Size = 0;
+ uint64_t ProcessID = 0;
+ StringRef FileName;
+ uri_type_t URIType;
+
+ // Constructors
+ // TODO: add a Copy ctor ?
+ OffloadBundleURI(StringRef file, int64_t off, int64_t size)
+ : Offset(off), Size(size), ProcessID(0), FileName(file),
+ URIType(FILE_URI) {}
+
+ OffloadBundleURI(StringRef str, uri_type_t type) {
+ URIType = type;
+ switch (URIType) {
+ case FILE_URI:
+ parseFileName(str);
+ break;
+ case MEMORY_URI:
+ parseMemoryURI(str);
+ break;
+ default:
+ report_fatal_error("Unrecognized URI type.");
+ }
+ }
+
+ void parseFileName(StringRef str) {
+ ProcessID = 0;
+ URIType = FILE_URI;
+ if (str.consume_front("file://")) {
+ StringRef FilePathname =
+ str.take_until([](char c) { return (c == '#') || (c == '?'); });
+ FileName = FilePathname;
+ str = str.drop_front(FilePathname.size());
+
+ if (str.consume_front("#offset=")) {
+ StringRef OffsetStr = str.take_until([](char c) { return c == '&'; });
+ OffsetStr.getAsInteger(10, Offset);
+ str = str.drop_front(OffsetStr.size());
+
+ if (str.consume_front("&size=")) {
+ Size;
+ str.getAsInteger(10, Size);
+ } else
+ report_fatal_error("Reading 'size' in URI.");
+ } else
+ report_fatal_error("Reading 'offset' in URI.");
+ } else
+ report_fatal_error("Reading type of URI.");
+ }
+
+ void parseMemoryURI(StringRef str) {
+ // TODO: add parseMemoryURI type
+ }
+
+ StringRef getFileName() const { return FileName; }
+};
+
/// Extracts embedded device offloading code from a memory \p Buffer to a list
/// of \p Binaries.
Error extractOffloadBinaries(MemoryBufferRef Buffer,
SmallVectorImpl<OffloadFile> &Binaries);
+Error extractFatBinaryFromObject(const ObjectFile &Obj,
+ SmallVectorImpl<OffloadFatBinBundle> &Bundles);
+
+Error extractCodeObject(const ObjectFile &Source, int64_t Offset, int64_t Size,
+ StringRef OutputFileName);
+
+Error extractURI(StringRef URIstr);
+
/// Convert a string \p Name to an image kind.
ImageKind getImageKind(StringRef Name);
diff --git a/llvm/lib/Object/OffloadBinary.cpp b/llvm/lib/Object/OffloadBinary.cpp
index 89dc12551494fd..5b30fcbdc19732 100644
--- a/llvm/lib/Object/OffloadBinary.cpp
+++ b/llvm/lib/Object/OffloadBinary.cpp
@@ -9,6 +9,7 @@
#include "llvm/Object/OffloadBinary.h"
#include "llvm/ADT/StringSwitch.h"
+#include "llvm/BinaryFormat/COFF.h"
#include "llvm/BinaryFormat/Magic.h"
#include "llvm/IR/Constants.h"
#include "llvm/IR/Module.h"
@@ -23,14 +24,20 @@
#include "llvm/Object/IRObjectFile.h"
#include "llvm/Object/ObjectFile.h"
#include "llvm/Support/Alignment.h"
+#include "llvm/Support/BinaryStreamReader.h"
#include "llvm/Support/FileOutputBuffer.h"
#include "llvm/Support/SourceMgr.h"
+#include "llvm/Support/Timer.h"
using namespace llvm;
using namespace llvm::object;
namespace {
+static llvm::TimerGroup
+ ClangOffloadBundlerTimerGroup("Clang Offload Bundler Timer Group",
+ "Timer group for clang offload bundler");
+
/// Attempts to extract all the embedded device images contained inside the
/// buffer \p Contents. The buffer is expected to contain a valid offloading
/// binary format.
@@ -99,6 +106,48 @@ Error extractFromObject(const ObjectFile &Obj,
return Error::success();
}
+// Extract an Offload bundle (usually a Clang Offload Bundle) from a fat_bin
+// section
+Error extractOffloadBundle(MemoryBufferRef Contents, uint64_t SectionOffset,
+ StringRef fileName,
+ SmallVectorImpl<OffloadFatBinBundle> &Bundles) {
+
+ uint64_t Offset = 0;
+ int64_t nextbundleStart = 0;
+
+ // There could be multiple offloading bundles stored at this section.
+ while (nextbundleStart >= 0) {
+
+ std::unique_ptr<MemoryBuffer> Buffer =
+ MemoryBuffer::getMemBuffer(Contents.getBuffer().drop_front(Offset), "",
+ /*RequiresNullTerminator*/ false);
+
+ // Create the FatBinBindle object. This will also create the Bundle Entry
+ // list info.
+ auto FatBundleOrErr =
+ OffloadFatBinBundle::create(*Buffer, SectionOffset + Offset, fileName);
+ if (!FatBundleOrErr)
+ return FatBundleOrErr.takeError();
+ OffloadFatBinBundle &Bundle = **FatBundleOrErr;
+
+ // add current Bundle to list.
+ Bundles.emplace_back(std::move(**FatBundleOrErr));
+
+ // find the next bundle by searching for the magic string
+ StringRef str = Buffer->getBuffer();
+ nextbundleStart =
+ (int64_t)str.find(StringRef("__CLANG_OFFLOAD_BUNDLE__"), 24);
+
+ if (nextbundleStart >= 0)
+ Offset += nextbundleStart;
+ else {
+ return Error::success();
+ }
+ } // end of while loop
+
+ return Error::success();
+}
+
Error extractFromBitcode(MemoryBufferRef Buffer,
SmallVectorImpl<OffloadFile> &Binaries) {
LLVMContext Context;
@@ -170,6 +219,103 @@ Error extractFromArchive(const Archive &Library,
} // namespace
+Error OffloadFatBinBundle::ReadEntries(StringRef Buffer,
+ uint64_t SectionOffset) {
+ uint64_t BundleNumber = 0;
+ uint64_t NumOfEntries = 0;
+
+ // get Reader
+ BinaryStreamReader Reader(Buffer, llvm::endianness::little);
+
+ // Read the Magic String first.
+ StringRef Magic;
+ if (auto EC = Reader.readFixedString(Magic, 24)) {
+ return errorCodeToError(object_error::parse_failed);
+ }
+
+ // read the number of Code Objects (Entries) in the current Bundle.
+ if (auto EC = Reader.readInteger(NumOfEntries)) {
+ printf("OffloadFatBinBundle::ReadEntries .... failed to read number of "
+ "Entries\n");
+ return errorCodeToError(object_error::parse_failed);
+ }
+ NumberOfEntries = NumOfEntries;
+
+ // For each Bundle Entry (code object)
+ for (uint64_t I = 0; I < NumOfEntries; I++) {
+ uint64_t EntrySize;
+ uint64_t EntryOffset;
+ uint64_t EntryIDSize;
+ StringRef EntryID;
+ uint64_t absOffset;
+
+ if (auto EC = Reader.readInteger(EntryOffset)) {
+ return errorCodeToError(object_error::parse_failed);
+ }
+
+ if (auto EC = Reader.readInteger(EntrySize)) {
+ return errorCodeToError(object_error::parse_failed);
+ }
+
+ if (auto EC = Reader.readInteger(EntryIDSize)) {
+ return errorCodeToError(object_error::parse_failed);
+ }
+
+ if (auto EC = Reader.readFixedString(EntryID, EntryIDSize)) {
+ return errorCodeToError(object_error::parse_failed);
+ }
+
+ // create a Bundle Entry object:
+ auto entry = new OffloadFatBinBundle::BundleEntry(
+ EntryOffset + SectionOffset, EntrySize, EntryIDSize, EntryID);
+
+ Entries->push_back(*entry);
+ } // end of for loop
+
+ return Error::success();
+}
+
+Expected<std::unique_ptr<OffloadFatBinBundle>>
+OffloadFatBinBundle::create(MemoryBufferRef Buf, uint64_t SectionOffset,
+ StringRef fileName) {
+ if (Buf.getBufferSize() < 24)
+ return errorCodeToError(object_error::parse_failed);
+
+ // Check for magic bytes.
+ if (identify_magic(Buf.getBuffer()) != file_magic::offload_bundle)
+ return errorCodeToError(object_error::parse_failed);
+
+ OffloadFatBinBundle *TheBundle = new OffloadFatBinBundle(Buf, fileName);
+
+ // Read the Bundle Entries
+ Error Err = TheBundle->ReadEntries(Buf.getBuffer(), SectionOffset);
+ if (Err)
+ return errorCodeToError(object_error::parse_failed);
+
+ return std::unique_ptr<OffloadFatBinBundle>(TheBundle);
+}
+
+Error OffloadFatBinBundle::extractBundle(const ObjectFile &Source) {
+ // This will extract all entries in the Bundle
+ SmallVectorImpl<OffloadFatBinBundle::BundleEntry>::iterator it =
+ Entries->begin();
+ for (int64_t I = 0; I < getNumEntries(); I++) {
+
+ if (it->Size > 0) {
+ // create output file name. Which should be
+ // <fileName>-offset<Offset>-size<Size>.co"
+ std::string str = getFileName().str() + "-offset" + itostr(it->Offset) +
+ "-size" + itostr(it->Size) + ".co";
+ if (Error Err = object::extractCodeObject(Source, it->Offset, it->Size,
+ StringRef(str)))
+ return Err;
+ }
+ ++it;
+ }
+
+ return Error::success();
+}
+
Expected<std::unique_ptr<OffloadBinary>>
OffloadBinary::create(MemoryBufferRef Buf) {
if (Buf.getBufferSize() < sizeof(Header) + sizeof(Entry))
@@ -299,6 +445,108 @@ Error object::extractOffloadBinaries(MemoryBufferRef Buffer,
}
}
+Error object::extractFatBinaryFromObject(
+ const ObjectFile &Obj, SmallVectorImpl<OffloadFatBinBundle> &Bundles) {
+ assert((Obj.isELF() || Obj.isCOFF()) && "Invalid file type");
+
+ // iterate through Sections until we find an offload_bundle section.
+ for (SectionRef Sec : Obj.sections()) {
+ Expected<StringRef> Buffer = Sec.getContents();
+ if (!Buffer)
+ return Buffer.takeError();
+
+ // If it does not start with the reserved suffix, just skip this section.
+ if ((llvm::identify_magic(*Buffer) == llvm::file_magic::offload_bundle) ||
+ (llvm::identify_magic(*Buffer) ==
+ llvm::file_magic::offload_bundle_compressed)) {
+
+ uint64_t SectionOffset = 0;
+ if (Obj.isELF()) {
+ SectionOffset = ELFSectionRef(Sec).getOffset();
+ } else if (Obj.isCOFF()) {
+ if (const COFFObjectFile *COFFObj = dyn_cast<COFFObjectFile>(&Obj)) {
+ const coff_section *CoffSection = COFFObj->getCOFFSection(Sec);
+ fprintf(
+ stderr, "DAVE: COFF viritual address =0x%llX\n",
+ CoffSection
+ ->VirtualAddress); // COFFObj->getCOFFSection(Sec)->VirtualAddress);
+ }
+ }
+
+ MemoryBufferRef Contents(*Buffer, Obj.getFileName());
+
+ if (llvm::identify_magic(*Buffer) ==
+ llvm::file_magic::offload_bundle_compressed) {
+ // Decompress the input if necessary.
+ Expected<std::unique_ptr<MemoryBuffer>> DecompressedBufferOrErr =
+ CompressedOffloadBundle::decompress(Contents, false);
+
+ if (!DecompressedBufferOrErr)
+ return createStringError(
+ inconvertibleErrorCode(),
+ "Failed to decompress input: " +
+ llvm::toString(DecompressedBufferOrErr.takeError()));
+
+ MemoryBuffer &DecompressedInput = **DecompressedBufferOrErr;
+ if (Error Err = extractOffloadBundle(DecompressedInput, SectionOffset,
+ Obj.getFileName(), Bundles))
+ return Err;
+ } else {
+ if (Error Err = extractOffloadBundle(Contents, SectionOffset,
+ Obj.getFileName(), Bundles))
+ return Err;
+ }
+ }
+ }
+ return Error::success();
+}
+
+Error object::extractCodeObject(const ObjectFile &Source, int64_t Offset,
+ int64_t Size, StringRef OutputFileName) {
+ Expected<std::unique_ptr<FileOutputBuffer>> BufferOrErr =
+ FileOutputBuffer::create(OutputFileName, Size);
+
+ if (!BufferOrErr)
+ return BufferOrErr.takeError();
+
+ Expected<MemoryBufferRef> InputBuffOrErr = Source.getMemoryBufferRef();
+ if (Error Err = InputBuffOrErr.takeError())
+ return Err;
+
+ std::unique_ptr<FileOutputBuffer> Buf = std::move(*BufferOrErr);
+ std::copy(InputBuffOrErr->getBufferStart() + Offset,
+ InputBuffOrErr->getBufferStart() + Offset + Size,
+ Buf->getBufferStart());
+ if (Error E = Buf->commit())
+ return E;
+
+ return Error::success();
+}
+
+// given a file name, offset, and size, extract data into a code object file,
+// into file <SourceFile>-offset<Offset>-size<Size>.co
+Error object::extractURI(StringRef URIstr) {
+ // create a URI object
+ object::OffloadBundleURI *uri =
+ new object::OffloadBundleURI(URIstr, FILE_URI);
+
+ std::string OutputFile = uri->FileName.str();
+ OutputFile +=
+ "-offset" + itostr(uri->Offset) + "-size" + itostr(uri->Size) + ".co";
+
+ // Create an ObjectFile object from uri.file_uri
+ auto ObjOrErr = ObjectFile::createObjectFile(uri->FileName);
+ if (!ObjOrErr)
+ return ObjOrErr.takeError();
+
+ auto Obj = ObjOrErr->getBinary();
+ if (Error Err =
+ object::extractCodeObject(*Obj, uri->Offset, uri->Size, OutputFile))
+ return Err;
+
+ return Error::success();
+}
+
OffloadKind object::getOffloadKind(StringRef Name) {
return llvm::StringSwitch<OffloadKind>(Name)
.Case("openmp", OFK_OpenMP)
@@ -382,3 +630,223 @@ bool object::areTargetsCompatible(const OffloadFile::TargetID &LHS,
return false;
return true;
}
+
+// Utility function to format numbers with commas
+static std::string formatWithCommas(unsigned long long Value) {
+ std::string Num = std::to_string(Value);
+ int InsertPosition = Num.length() - 3;
+ while (InsertPosition > 0) {
+ Num.insert(InsertPosition, ",");
+ InsertPosition -= 3;
+ }
+ return Num;
+}
+
+llvm::Expected<std::unique_ptr<llvm::MemoryBuffer>>
+CompressedOffloadBundle::decompress(llvm::MemoryBufferRef &Input,
+
+ bool Verbose) {
+ StringRef Blob = Input.getBuffer();
+
+ if (Blob.size() < V1HeaderSize)
+ return llvm::MemoryBuffer::getMemBufferCopy(Blob);
+
+ if (llvm::identify_magic(Blob) !=
+ llvm::file_magic::offload_bundle_compressed) {
+ if (Verbose)
+ llvm::errs() << "Uncompressed bundle.\n";
+ return llvm::MemoryBuffer::getMemBufferCopy(Blob);
+ }
+
+ size_t CurrentOffset = MagicSize;
+
+ uint16_t ThisVersion;
+ memcpy(&ThisVersion, Blob.data() + CurrentOffset, sizeof(uint16_t));
+ CurrentOffset += VersionFieldSize;
+
+ uint16_t CompressionMethod;
+ memcpy(&CompressionMethod, Blob.data() + CurrentOffset, sizeof(uint16_t));
+ CurrentOffset += MethodFieldSize;
+
+ uint32_t TotalFileSize;
+ if (ThisVersion >= 2) {
+ if (Blob.size() < V2HeaderSize)
+ return createStringError(inconvertibleErrorCode(),
+ "Compressed bundle header size too small");
+ memcpy(&TotalFileSize, Blob.data() + CurrentOffset, sizeof(uint32_t));
+ CurrentOffset += FileSizeFieldSize;
+ }
+
+ uint32_t UncompressedSize;
+ memcpy(&UncompressedSize, Blob.data() + CurrentOffset, sizeof(uint32_t));
+ CurrentOffset += UncompressedSizeFieldSize;
+
+ uint64_t StoredHash;
+ memcpy(&StoredHash, Blob.data() + CurrentOffset, sizeof(uint64_t));
+ CurrentOffset += HashFieldSize;
+
+ llvm::compression::Format CompressionFormat;
+ if (CompressionMethod ==
+ static_cast<uint16_t>(llvm::compression::Format::Zlib))
+ CompressionFormat = llvm::compression::Format::Zlib;
+ else if (CompressionMethod ==
+ static_cast<uint16_t>(llvm::compression::Format::Zstd))
+ CompressionFormat = llvm::compression::Format::Zstd;
+ else
+ return createStringError(inconvertibleErrorCode(),
+ "Unknown compressing method");
+
+ llvm::Timer DecompressTimer("Decompression Timer", "Decompression time",
+ ClangOffloadBundlerTimerGroup);
+ if (Verbose)
+ DecompressTimer.startTimer();
+
+ SmallVector<uint8_t, 0> DecompressedData;
+ StringRef CompressedData = Blob.substr(CurrentOffset);
+ if (llvm::Error DecompressionError = llvm::compression::decompress(
+ CompressionFormat, llvm::arrayRefFromStringRef(CompressedData),
+ DecompressedData, UncompressedSize))
+ return createStringError(inconvertibleErrorCode(),
+ "Could not decompress embedded file contents: " +
+ llvm::toString(std::move(DecompressionError)));
+
+ if (Verbose) {
+ DecompressTimer.stopTimer();
+
+ double DecompressionTimeSeconds =
+ DecompressTimer.getTotalTime().getWallTime();
+
+ // Recalculate MD5 hash for integrity check
+ llvm::Timer HashRecalcTimer("Hash Recalculation Timer",
+ "Hash recalculation time",
+ ClangOffloadBundlerTimerGroup);
+ HashRecalcTimer.startTimer();
+ llvm::MD5 Hash;
+ llvm::MD5::MD5Result Result;
+ Hash.update(llvm::ArrayRef<uint8_t>(DecompressedData.data(),
+ DecompressedData.size()));
+ Hash.final(Result);
+ uint64_t RecalculatedHash = Result.low();
+ HashRecalcTimer.stopTimer();
+ bool HashMatch = (StoredHash == RecalculatedHash);
+
+ double CompressionRate =
+ static_cast<double>(UncompressedSize) / CompressedData.size();
+ double DecompressionSpeedMBs =
+ (UncompressedSize / (1024.0 * 1024.0)) / DecompressionTimeSeconds;
+
+ llvm::errs() << "Compressed bundle format version: " << ThisVersion << "\n";
+ if (ThisVersion >= 2)
+ llvm::errs() << "Total file size (from header): "
+ << formatWithCommas(TotalFileSize) << " bytes\n";
+ llvm::errs() << "Decompression method: "
+ << (CompressionFormat == llvm::compression::Format::Zlib
+ ? "zlib"
+ : "zstd")
+ << "\n"
+ << "Size before decompression: "
+ << formatWithCommas(CompressedData.size()) << " bytes\n"
+ << "Size after decompression: "
+ << formatWithCommas(UncompressedSize) << " bytes\n"
+ << "Compression rate: "
+ << llvm::format("%.2lf", CompressionRate) << "\n"
+ << "Compression ratio: "
+ << llvm::format("%.2lf%%", 100.0 / CompressionRate) << "\n"
+ << "Decompression speed: "
+ << llvm::format("%.2lf MB/s", DecompressionSpeedMBs) << "\n"
+ << "Stored hash: " << llvm::format_hex(StoredHash, 16) << "\n"
+ << "Recalculated hash: "
+ << llvm::format_hex(RecalculatedHash, 16) << "\n"
+ << "Hashes match: " << (HashMatch ? "Yes" : "No") << "\n";
+ }
+
+ return llvm::MemoryBuffer::getMemBufferCopy(
+ llvm::toStringRef(DecompressedData));
+}
+
+llvm::Expected<std::unique_ptr<llvm::MemoryBuffer>>
+CompressedOffloadBundle::compress(llvm::compression::Params P,
+ const llvm::MemoryBuffer &Input,
+ bool Verbose) {
+ if (!llvm::compression::zstd::isAvailable() &&
+ !llvm::compression::zlib::isAvailable())
+ return createStringError(llvm::inconvertibleErrorCode(),
+ "Compression not supported");
+
+ llvm::Timer HashTimer("Hash Calculation Timer", "Hash calculation time",
+ ClangOffloadBundlerTimerGroup);
+ if (Verbose)
+ HashTimer.startTimer();
+ llvm::MD5 Hash;
+ llvm::MD5::MD5Result Result;
+ Hash.update(Input.getBuffer());
+ Hash.final(Result);
+ uint64_t TruncatedHash = Result.low();
+ if (Verbose)
+ HashTimer.stopTimer();
+
+ SmallVector<uint8_t, 0> CompressedBuffer;
+ auto BufferUint8 = llvm::ArrayRef<uint8_t>(
+ reinterpret_cast<const uint8_t *>(Input.getBuffer().data()),
+ Input.getBuffer().size());
+
+ llvm::Timer CompressTimer("Compression Timer", "Compression time",
+ ClangOffloadBundlerTimerGroup);
+ if (Verbose)
+ CompressTimer.startTimer();
+ llvm::compression::compress(P, BufferUint8, CompressedBuffer);
+ if (Verbose)
+ CompressTimer.stopTimer();
+
+ uint16_t CompressionMethod = static_cast<uint16_t>(P.format);
+ uint32_t UncompressedSize = Input.getBuffer().size();
+ uint32_t TotalFileSize = MagicNumber.size() + sizeof(TotalFileSize) +
+ sizeof(Version) + sizeof(CompressionMethod) +
+ sizeof(UncompressedSize) + sizeof(TruncatedHash) +
+ CompressedBuffer.size();
+
+ SmallVector<char, 0> FinalBuffer;
+ llvm::raw_svector_ostream OS(FinalBuffer);
+ OS << MagicNumber;
+ OS.write(reinterpret_cast<const char *>(&Version), sizeof(Version));
+ OS.write(reinterpret_cast<const char *>(&CompressionMethod),
+ sizeof(CompressionMethod));
+ OS.write(reinterpret_cast<const char *>(&TotalFileSize),
+ sizeof(TotalFileSize));
+ OS.write(reinterpret_cast<const char *>(&UncompressedSize),
+ sizeof(UncompressedSize));
+ OS.write(reinterpret_cast<const char *>(&TruncatedHash),
+ sizeof(TruncatedHash));
+ OS.write(reinterpret_cast<const char *>(CompressedBuffer.data()),
+ CompressedBuffer.size());
+
+ if (Verbose) {
+ auto MethodUsed =
+ P.format == llvm::compression::Format::Zstd ? "zstd" : "zlib";
+ double CompressionRate =
+ static_cast<double>(UncompressedSize) / CompressedBuffer.size();
+ double CompressionTimeSeconds = CompressTimer.getTotalTime().getWallTime();
+ double CompressionSpeedMBs =
+ (UncompressedSize / (1024.0 * 1024.0)) / CompressionTimeSeconds;
+
+ llvm::errs() << "Compressed bundle format version: " << Version << "\n"
+ << "Total file size (including headers): "
+ << formatWithCommas(TotalFileSize) << " bytes\n"
+ << "Compression method used: " << MethodUsed << "\n"
+ << "Compression level: " << P.level << "\n"
+ << "Binary size before compression: "
+ << formatWithCommas(UncompressedSize) << " bytes\n"
+ << "Binary size after compression: "
+ << formatWithCommas(CompressedBuffer.size()) << " bytes\n"
+ << "Compression rate: "
+ << llvm::format("%.2lf", CompressionRate) << "\n"
+ << "Compression ratio: "
+ << llvm::format("%.2lf%%", 100.0 / CompressionRate) << "\n"
+ << "Compression speed: "
+ << llvm::format("%.2lf MB/s", CompressionSpeedMBs) << "\n"
+ << "Truncated MD5 hash: "
+ << llvm::format_hex(TruncatedHash, 16) << "\n";
+ }
+ return llvm::MemoryBuffer::getMemBufferCopy(
+ llvm::StringRef(FinalBuffer.data(), FinalBuffer.size()));
+}
>From eea2802e312e73160efacc6e4d6c048806de83f2 Mon Sep 17 00:00:00 2001
From: David <dsalinas at amd.com>
Date: Wed, 28 Aug 2024 04:01:38 +0100
Subject: [PATCH 2/3] Extend LLVM Offloading API for binary fatbin Bundles
With the intention to provide a common API for offloading, this
extension to the existing LLVM Offloading API adds support for
Binary Fatbin Bundles; moving some support from the Clang offloading
API. The intention is to add functionality to LLVM tooling for
Binary Fatbin Bundles in subsequent commits.
Change-Id: I907fdcbcd0545162a0ce1cf17ebf7c9f3a4dbde6
---
llvm/lib/Object/OffloadBinary.cpp | 1 +
1 file changed, 1 insertion(+)
diff --git a/llvm/lib/Object/OffloadBinary.cpp b/llvm/lib/Object/OffloadBinary.cpp
index 5b30fcbdc19732..cd7cf2504fd056 100644
--- a/llvm/lib/Object/OffloadBinary.cpp
+++ b/llvm/lib/Object/OffloadBinary.cpp
@@ -18,6 +18,7 @@
#include "llvm/Object/Archive.h"
#include "llvm/Object/ArchiveWriter.h"
#include "llvm/Object/Binary.h"
+#include "llvm/BinaryFormat/COFF.h"
#include "llvm/Object/COFF.h"
#include "llvm/Object/ELFObjectFile.h"
#include "llvm/Object/Error.h"
>From 633c8af53442a365a98d1cda2405fb2dc098592f Mon Sep 17 00:00:00 2001
From: dsalinas <dsalinas at MKM-L1-DSALINAS.amd.com>
Date: Mon, 23 Sep 2024 14:31:50 -0400
Subject: [PATCH 3/3] Extend llvm-objdump to support FatBins
add option --offload-fatbin
SWDEV-333176 - Shift functionality of 'roc-obj-*' perl scripts into
llvm-objdump
Change-Id: Ibc865f80e30aa1a6e5495ecfe617be68a5e15fcf
---
llvm/include/llvm/Object/OffloadBinary.h | 62 +-
llvm/lib/Object/ObjectFile.cpp | 1 -
llvm/lib/Object/OffloadBinary.cpp | 13 +-
.../Offloading/fatbin-offloading.test | 80 ++
.../tools/llvm-objdump/Offloading/fatbin.test | 844 ++++++++++++++++++
llvm/tools/llvm-objdump/ObjdumpOpts.td | 3 +
llvm/tools/llvm-objdump/OffloadDump.cpp | 52 ++
llvm/tools/llvm-objdump/OffloadDump.h | 3 +-
llvm/tools/llvm-objdump/llvm-objdump.cpp | 7 +-
9 files changed, 1030 insertions(+), 35 deletions(-)
create mode 100644 llvm/test/tools/llvm-objdump/Offloading/fatbin-offloading.test
create mode 100644 llvm/test/tools/llvm-objdump/Offloading/fatbin.test
diff --git a/llvm/include/llvm/Object/OffloadBinary.h b/llvm/include/llvm/Object/OffloadBinary.h
index ee01cabe146517..797dfc71e71c99 100644
--- a/llvm/include/llvm/Object/OffloadBinary.h
+++ b/llvm/include/llvm/Object/OffloadBinary.h
@@ -210,51 +210,54 @@ class OffloadFile : public OwningBinary<OffloadBinary> {
}
};
+struct BundleEntry {
+ uint64_t Offset = 0u;
+ uint64_t Size = 0u;
+ uint64_t IDLength = 0u;
+ StringRef ID;
+ BundleEntry(uint64_t O, uint64_t S, uint64_t I, StringRef T)
+ : Offset(O), Size(S), IDLength(I), ID(T) {}
+ void dumpInfo(raw_ostream &OS) {
+ OS << "Offset = " << Offset << ", Size = " << Size
+ << ", ID Length = " << IDLength << ", ID = " << ID;
+ }
+ void dumpURI(raw_ostream &OS, StringRef filePath) {
+ OS << ID.data() << "\tfile:\/\/" << filePath << "#offset=" << Offset
+ << "&size=" << Size << "\n";
+ }
+};
+
class OffloadFatBinBundle {
private:
uint64_t Size = 0u;
StringRef FileName;
int64_t NumberOfEntries;
+ SmallVector<BundleEntry> Entries;
public:
- struct BundleEntry {
- uint64_t Offset = 0u;
- uint64_t Size = 0u;
- uint64_t IDLength = 0u;
- StringRef ID;
- BundleEntry(uint64_t O, uint64_t S, uint64_t I, StringRef T)
- : Offset(O), Size(S), IDLength(I), ID(T) {}
- void dump(raw_ostream &OS) {
- OS << "Offset = " << Offset << ", Size = " << Size
- << ", ID Length = " << IDLength << ", ID = " << ID;
- }
- void dumpURI(raw_ostream &OS, StringRef filePath) {
- OS << ID.data() << "\tfile:\/\/" << filePath << "#offset=" << Offset
- << "&size=" << Size << "\n";
- }
- };
-
+ SmallVector<BundleEntry> getEntries() { return Entries; }
uint64_t getSize() const { return Size; }
StringRef getFileName() const { return FileName; }
int64_t getNumEntries() const { return NumberOfEntries; }
- std::unique_ptr<SmallVector<BundleEntry>> Entries;
static Expected<std::unique_ptr<OffloadFatBinBundle>>
create(MemoryBufferRef, uint64_t SectionOffset, StringRef fileName);
Error extractBundle(const ObjectFile &Source);
+ Error DumpEntryToCodeObject();
+
Error ReadEntries(StringRef Section, uint64_t SectionOffset);
void DumpEntries() {
- SmallVectorImpl<BundleEntry>::iterator it = Entries->begin();
- for (int64_t I = 0; I < Entries->size(); I++) {
- it->dump(outs());
+ SmallVectorImpl<BundleEntry>::iterator it = Entries.begin();
+ for (int64_t I = 0; I < Entries.size(); I++) {
+ it->dumpInfo(outs());
++it;
}
}
void PrintEntriesAsURI() {
- SmallVectorImpl<BundleEntry>::iterator it = Entries->begin();
+ SmallVectorImpl<BundleEntry>::iterator it = Entries.begin();
for (int64_t I = 0; I < NumberOfEntries; I++) {
it->dumpURI(outs(), FileName);
++it;
@@ -263,7 +266,20 @@ class OffloadFatBinBundle {
OffloadFatBinBundle(MemoryBufferRef Source, StringRef file) : FileName(file) {
NumberOfEntries = 0;
- Entries = std::make_unique<SmallVector<BundleEntry>>();
+ Entries = SmallVector<BundleEntry>();
+ }
+
+ SmallVector<BundleEntry> EntryIDContains(StringRef str) {
+ SmallVector<BundleEntry> found = SmallVector<BundleEntry>();
+ SmallVectorImpl<BundleEntry>::iterator it = Entries.begin();
+ for (int64_t I = 0; I < NumberOfEntries; I++) {
+ if (it->ID.contains(str)) {
+ found.push_back(*it);
+ }
+
+ ++it;
+ }
+ return found;
}
};
diff --git a/llvm/lib/Object/ObjectFile.cpp b/llvm/lib/Object/ObjectFile.cpp
index 6a226a3bbdbca3..636e3e2423d53f 100644
--- a/llvm/lib/Object/ObjectFile.cpp
+++ b/llvm/lib/Object/ObjectFile.cpp
@@ -212,7 +212,6 @@ ObjectFile::createObjectFile(StringRef ObjectPath) {
if (std::error_code EC = FileOrErr.getError())
return errorCodeToError(EC);
std::unique_ptr<MemoryBuffer> Buffer = std::move(FileOrErr.get());
-
Expected<std::unique_ptr<ObjectFile>> ObjOrErr =
createObjectFile(Buffer->getMemBufferRef());
if (Error Err = ObjOrErr.takeError())
diff --git a/llvm/lib/Object/OffloadBinary.cpp b/llvm/lib/Object/OffloadBinary.cpp
index cd7cf2504fd056..0ed0bf95b926c2 100644
--- a/llvm/lib/Object/OffloadBinary.cpp
+++ b/llvm/lib/Object/OffloadBinary.cpp
@@ -267,10 +267,10 @@ Error OffloadFatBinBundle::ReadEntries(StringRef Buffer,
}
// create a Bundle Entry object:
- auto entry = new OffloadFatBinBundle::BundleEntry(
- EntryOffset + SectionOffset, EntrySize, EntryIDSize, EntryID);
+ auto entry = new BundleEntry(EntryOffset + SectionOffset, EntrySize,
+ EntryIDSize, EntryID);
- Entries->push_back(*entry);
+ Entries.push_back(*entry);
} // end of for loop
return Error::success();
@@ -298,8 +298,7 @@ OffloadFatBinBundle::create(MemoryBufferRef Buf, uint64_t SectionOffset,
Error OffloadFatBinBundle::extractBundle(const ObjectFile &Source) {
// This will extract all entries in the Bundle
- SmallVectorImpl<OffloadFatBinBundle::BundleEntry>::iterator it =
- Entries->begin();
+ SmallVectorImpl<BundleEntry>::iterator it = Entries.begin();
for (int64_t I = 0; I < getNumEntries(); I++) {
if (it->Size > 0) {
@@ -467,10 +466,6 @@ Error object::extractFatBinaryFromObject(
} else if (Obj.isCOFF()) {
if (const COFFObjectFile *COFFObj = dyn_cast<COFFObjectFile>(&Obj)) {
const coff_section *CoffSection = COFFObj->getCOFFSection(Sec);
- fprintf(
- stderr, "DAVE: COFF viritual address =0x%llX\n",
- CoffSection
- ->VirtualAddress); // COFFObj->getCOFFSection(Sec)->VirtualAddress);
}
}
diff --git a/llvm/test/tools/llvm-objdump/Offloading/fatbin-offloading.test b/llvm/test/tools/llvm-objdump/Offloading/fatbin-offloading.test
new file mode 100644
index 00000000000000..195bf38d985bc0
--- /dev/null
+++ b/llvm/test/tools/llvm-objdump/Offloading/fatbin-offloading.test
@@ -0,0 +1,80 @@
+// RUN: clang++ -x hip --hip-link --offload-arch=gfx1100 --offload-arch=gfx1101 -o %t %s
+// RUN: llvm-objdump %t --offload-fatbin
+// RUN: llvm-objdump %t --offload-fatbin --arch-name=gfx1100
+
+#include <stdio.h>
+
+#include <iostream>
+#include "hip/hip_runtime.h"
+
+__global__ void simpleAdd(uint32_t* A_d, const uint32_t* B_d, size_t N)
+{
+ size_t i = (blockIdx.x * blockDim.x + threadIdx.x);
+ A_d[i] += B_d[i];
+}
+
+int main()
+{
+ int device_count = 0;
+ hipGetDeviceCount(&device_count);
+
+ std::cout<< "Found " << device_count << " HIP devices." << std::endl;
+ for( int i = 0; i < device_count; i++ )
+ {
+ std::cout << "Device " << i << std::endl;
+ hipDeviceProp_t props;
+ hipGetDeviceProperties(&props, i);
+ std::cout << " Name: " << props.name << std::endl;
+ }
+
+ hipSetDevice(0);
+
+ uint32_t *A_d, *B_d;
+ uint32_t *A_h, *B_h;
+
+ size_t N = 1000;
+ size_t Nbytes = N * sizeof( uint32_t );
+
+ // Allocating host memory
+ A_h = (uint32_t*) malloc(Nbytes);
+ B_h = (uint32_t*) malloc(Nbytes);
+
+ for( size_t i = 0; i < N ; i++ )
+ {
+ A_h[i] = i;
+ B_h[i] = 2*i;
+ }
+
+ // Allocating device memory
+ hipMalloc(&A_d, Nbytes);
+ hipMalloc(&B_d, Nbytes);
+
+ // Copy host to device
+ hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice);
+ hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice);
+
+ // launch kernel
+ const unsigned blocks = 512;
+ const unsigned threadsPerBlock = 256;
+ hipLaunchKernelGGL(simpleAdd, dim3(blocks), dim3(threadsPerBlock), 0, 0, A_d, B_d, N);
+
+ // Copy device to host
+ hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost);
+
+ // Verify
+ for( size_t i = 0; i < N; i++ )
+ {
+ std::cout << A_h[i] << " ";
+ uint32_t A_ref = 3*i;
+ if( A_h[i] != A_ref )
+ {
+ std::cout<< "Mismatch occured at " << i << ": " << A_h[i] << " != " << A_ref << std::endl;
+ break;
+ }
+ }
+ std::cout << std::endl;
+
+ // free up host memory
+ free( A_h );
+ free( B_h );
+}
diff --git a/llvm/test/tools/llvm-objdump/Offloading/fatbin.test b/llvm/test/tools/llvm-objdump/Offloading/fatbin.test
new file mode 100644
index 00000000000000..9e6a582bad9685
--- /dev/null
+++ b/llvm/test/tools/llvm-objdump/Offloading/fatbin.test
@@ -0,0 +1,844 @@
+## Test that --offload-fatbin works correctly
+
+# RUN: yaml2obj %s -o %t.elf
+# RUN: llvm-objdump --offload-fatbin %t.elf
+# RUN: llvm-objdump -d %t.elf:0.hipv4-amdgcn-amd-amdhsa--gfx908 | FileCheck %s
+
+
+# CHECK: s_load_dword s7, s[4:5], 0x24
+# CHECK-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0x0
+# CHECK-NEXT: v_mov_b32_e32 v1, 0
+# CHECK-NEXT: s_waitcnt lgkmcnt(0)
+# CHECK-NEXT: s_and_b32 s4, s7, 0xffff
+# CHECK-NEXT: s_mul_i32 s6, s6, s4
+# CHECK-NEXT: v_add_u32_e32 v0, s6, v0
+# CHECK-NEXT: v_lshlrev_b64 v[0:1], 2, v[0:1]
+# CHECK-NEXT: v_mov_b32_e32 v3, s3
+# CHECK-NEXT: v_add_co_u32_e32 v2, vcc, s2, v0
+# CHECK-NEXT: v_addc_co_u32_e32 v3, vcc, v3, v1, vcc
+# CHECK-NEXT: global_load_dword v2, v[2:3], off
+# CHECK-NEXT: v_mov_b32_e32 v3, s1
+# CHECK-NEXT: v_add_co_u32_e32 v0, vcc, s0, v0
+# CHECK-NEXT: v_addc_co_u32_e32 v1, vcc, v3, v1, vcc
+# CHECK-NEXT: global_load_dword v3, v[0:1], off
+# CHECK-NEXT: s_waitcnt vmcnt(0)
+# CHECK-NEXT: v_add_u32_e32 v2, v3, v2
+# CHECK-NEXT: global_store_dword v[0:1], v2, off
+# CHECK-NEXT: s_endpgm
+
+--- !ELF
+FileHeader:
+ Class: ELFCLASS64
+ Data: ELFDATA2LSB
+ Type: ET_EXEC
+ Machine: EM_X86_64
+ Entry: 0x2041B0
+ProgramHeaders:
+ - Type: PT_PHDR
+ Flags: [ PF_R ]
+ VAddr: 0x200040
+ Align: 0x8
+ Offset: 0x40
+ - Type: PT_INTERP
+ Flags: [ PF_R ]
+ FirstSec: .interp
+ LastSec: .interp
+ VAddr: 0x2002A8
+ Offset: 0x2A8
+ - Type: PT_LOAD
+ Flags: [ PF_R ]
+ FirstSec: .interp
+ LastSec: .eh_frame
+ VAddr: 0x200000
+ Align: 0x1000
+ Offset: 0x0
+ - Type: PT_LOAD
+ Flags: [ PF_X, PF_R ]
+ FirstSec: .text
+ LastSec: .plt
+ VAddr: 0x2041B0
+ Align: 0x1000
+ Offset: 0x31B0
+ - Type: PT_LOAD
+ Flags: [ PF_W, PF_R ]
+ FirstSec: .init_array
+ LastSec: .relro_padding
+ VAddr: 0x205AE0
+ Align: 0x1000
+ Offset: 0x3AE0
+ - Type: PT_LOAD
+ Flags: [ PF_W, PF_R ]
+ FirstSec: .data
+ LastSec: .bss
+ VAddr: 0x206CF8
+ Align: 0x1000
+ Offset: 0x3CF8
+ - Type: PT_DYNAMIC
+ Flags: [ PF_W, PF_R ]
+ FirstSec: .dynamic
+ LastSec: .dynamic
+ VAddr: 0x205B00
+ Align: 0x8
+ Offset: 0x3B00
+ - Type: PT_GNU_RELRO
+ Flags: [ PF_R ]
+ FirstSec: .init_array
+ LastSec: .relro_padding
+ VAddr: 0x205AE0
+ Offset: 0x3AE0
+ - Type: PT_GNU_EH_FRAME
+ Flags: [ PF_R ]
+ FirstSec: .eh_frame_hdr
+ LastSec: .eh_frame_hdr
+ VAddr: 0x202FE8
+ Align: 0x4
+ Offset: 0x2FE8
+ - Type: PT_GNU_STACK
+ Flags: [ PF_W, PF_R ]
+ Align: 0x0
+ Offset: 0x0
+ - Type: PT_NOTE
+ Flags: [ PF_R ]
+ FirstSec: .note.ABI-tag
+ LastSec: .note.ABI-tag
+ VAddr: 0x2002C4
+ Align: 0x4
+ Offset: 0x2C4
+Sections:
+ - Name: .interp
+ Type: SHT_PROGBITS
+ Flags: [ SHF_ALLOC ]
+ Address: 0x2002A8
+ AddressAlign: 0x1
+ Content: 2F6C696236342F6C642D6C696E75782D7838362D36342E736F2E3200
+ - Name: .note.ABI-tag
+ Type: SHT_NOTE
+ Flags: [ SHF_ALLOC ]
+ Address: 0x2002C4
+ AddressAlign: 0x4
+ Notes:
+ - Name: GNU
+ Desc: '00000000030000000200000000000000'
+ Type: NT_VERSION
+ - Name: .dynsym
+ Type: SHT_DYNSYM
+ Flags: [ SHF_ALLOC ]
+ Address: 0x2002E8
+ Link: .dynstr
+ AddressAlign: 0x8
+ - Name: .gnu.version
+ Type: SHT_GNU_versym
+ Flags: [ SHF_ALLOC ]
+ Address: 0x2005D0
+ Link: .dynsym
+ AddressAlign: 0x2
+ Entries: [ 0, 2, 1, 3, 4, 3, 5, 5, 5, 7, 6, 8, 6, 6, 9, 4,
+ 5, 4, 5, 5, 5, 7, 4, 6, 6, 4, 5, 5, 5, 6, 6 ]
+ - Name: .gnu.version_r
+ Type: SHT_GNU_verneed
+ Flags: [ SHF_ALLOC ]
+ Address: 0x200610
+ Link: .dynstr
+ AddressAlign: 0x4
+ Dependencies:
+ - Version: 1
+ File: libamdhip64.so.6
+ Entries:
+ - Name: hip_4.2
+ Hash: 252061554
+ Flags: 0
+ Other: 5
+ - Name: hip_6.0
+ Hash: 252062064
+ Flags: 0
+ Other: 9
+ - Version: 1
+ File: 'libstdc++.so.6'
+ Entries:
+ - Name: GLIBCXX_3.4
+ Hash: 143796596
+ Flags: 0
+ Other: 6
+ - Name: GLIBCXX_3.4.9
+ Hash: 36274057
+ Flags: 0
+ Other: 7
+ - Name: GLIBCXX_3.4.11
+ Hash: 43513953
+ Flags: 0
+ Other: 8
+ - Version: 1
+ File: libgcc_s.so.1
+ Entries:
+ - Name: GCC_3.0
+ Hash: 192489040
+ Flags: 0
+ Other: 3
+ - Version: 1
+ File: libc.so.6
+ Entries:
+ - Name: GLIBC_2.2.5
+ Hash: 157882997
+ Flags: 0
+ Other: 4
+ - Name: GLIBC_2.34
+ Hash: 110530996
+ Flags: 0
+ Other: 2
+ - Name: .gnu.hash
+ Type: SHT_GNU_HASH
+ Flags: [ SHF_ALLOC ]
+ Address: 0x2006D0
+ Link: .dynsym
+ AddressAlign: 0x8
+ Header:
+ SymNdx: 0x1D
+ Shift2: 0x1A
+ BloomFilter: [ 0x10000190000 ]
+ HashBuckets: [ 0x1D ]
+ HashValues: [ 0x430C9814, 0x4CD54529 ]
+ - Name: .dynstr
+ Type: SHT_STRTAB
+ Flags: [ SHF_ALLOC ]
+ Address: 0x2006F4
+ AddressAlign: 0x1
+ - Name: .rela.dyn
+ Type: SHT_RELA
+ Flags: [ SHF_ALLOC ]
+ Address: 0x200A18
+ Link: .dynsym
+ AddressAlign: 0x8
+ Relocations:
+ - Offset: 0x205CD0
+ Symbol: __libc_start_main
+ Type: R_X86_64_GLOB_DAT
+ - Offset: 0x205CD8
+ Symbol: __gmon_start__
+ Type: R_X86_64_GLOB_DAT
+ - Offset: 0x205CE0
+ Symbol: __register_frame_info
+ Type: R_X86_64_GLOB_DAT
+ - Offset: 0x205CE8
+ Symbol: __cxa_finalize
+ Type: R_X86_64_GLOB_DAT
+ - Offset: 0x205CF0
+ Symbol: __deregister_frame_info
+ Type: R_X86_64_GLOB_DAT
+ - Offset: 0x206E80
+ Symbol: _ZSt4cout
+ Type: R_X86_64_COPY
+ - Name: .rela.plt
+ Type: SHT_RELA
+ Flags: [ SHF_ALLOC, SHF_INFO_LINK ]
+ Address: 0x200AA8
+ Link: .dynsym
+ AddressAlign: 0x8
+ Info: .got.plt
+ Relocations:
+ - Offset: 0x206D20
+ Symbol: __register_frame_info
+ Type: R_X86_64_JUMP_SLOT
+ - Offset: 0x206D28
+ Symbol: __cxa_finalize
+ Type: R_X86_64_JUMP_SLOT
+ - Offset: 0x206D30
+ Symbol: __deregister_frame_info
+ Type: R_X86_64_JUMP_SLOT
+ - Offset: 0x206D38
+ Symbol: __hipPopCallConfiguration
+ Type: R_X86_64_JUMP_SLOT
+ - Offset: 0x206D40
+ Symbol: hipLaunchKernel
+ Type: R_X86_64_JUMP_SLOT
+ - Offset: 0x206D48
+ Symbol: hipGetDeviceCount
+ Type: R_X86_64_JUMP_SLOT
+ - Offset: 0x206D50
+ Symbol: _ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
+ Type: R_X86_64_JUMP_SLOT
+ - Offset: 0x206D58
+ Symbol: _ZNSolsEi
+ Type: R_X86_64_JUMP_SLOT
+ - Offset: 0x206D60
+ Symbol: _ZNKSt5ctypeIcE13_M_widen_initEv
+ Type: R_X86_64_JUMP_SLOT
+ - Offset: 0x206D68
+ Symbol: _ZNSo3putEc
+ Type: R_X86_64_JUMP_SLOT
+ - Offset: 0x206D70
+ Symbol: _ZNSo5flushEv
+ Type: R_X86_64_JUMP_SLOT
+ - Offset: 0x206D78
+ Symbol: hipGetDevicePropertiesR0600
+ Type: R_X86_64_JUMP_SLOT
+ - Offset: 0x206D80
+ Symbol: strlen
+ Type: R_X86_64_JUMP_SLOT
+ - Offset: 0x206D88
+ Symbol: hipSetDevice
+ Type: R_X86_64_JUMP_SLOT
+ - Offset: 0x206D90
+ Symbol: malloc
+ Type: R_X86_64_JUMP_SLOT
+ - Offset: 0x206D98
+ Symbol: hipMalloc
+ Type: R_X86_64_JUMP_SLOT
+ - Offset: 0x206DA0
+ Symbol: hipMemcpy
+ Type: R_X86_64_JUMP_SLOT
+ - Offset: 0x206DA8
+ Symbol: __hipPushCallConfiguration
+ Type: R_X86_64_JUMP_SLOT
+ - Offset: 0x206DB0
+ Symbol: _ZNSo9_M_insertImEERSoT_
+ Type: R_X86_64_JUMP_SLOT
+ - Offset: 0x206DB8
+ Symbol: free
+ Type: R_X86_64_JUMP_SLOT
+ - Offset: 0x206DC0
+ Symbol: _ZSt16__throw_bad_castv
+ Type: R_X86_64_JUMP_SLOT
+ - Offset: 0x206DC8
+ Symbol: _ZNSt8ios_base4InitC1Ev
+ Type: R_X86_64_JUMP_SLOT
+ - Offset: 0x206DD0
+ Symbol: _ZNSt8ios_base4InitD1Ev
+ Type: R_X86_64_JUMP_SLOT
+ - Offset: 0x206DD8
+ Symbol: __cxa_atexit
+ Type: R_X86_64_JUMP_SLOT
+ - Offset: 0x206DE0
+ Symbol: __hipRegisterFatBinary
+ Type: R_X86_64_JUMP_SLOT
+ - Offset: 0x206DE8
+ Symbol: __hipRegisterFunction
+ Type: R_X86_64_JUMP_SLOT
+ - Offset: 0x206DF0
+ Symbol: __hipUnregisterFatBinary
+ Type: R_X86_64_JUMP_SLOT
+ - Name: .rodata
+ Type: SHT_PROGBITS
+ Flags: [ SHF_ALLOC, SHF_MERGE, SHF_STRINGS ]
+ Address: 0x200D30
+ AddressAlign: 0x10
+ Content: 010002000000000000000000000000000800000008000000080000000800000000000000010000000200000003000000040000000400000004000000040000000C0000000C0000000C0000000C000000100000001000000010000000100000008042200000000000446576696365200020002048495020646576696365732E004D69736D61746368206F636375726564206174200020213D2000466F756E642000204E616D653A20005F5A3973696D706C65416464506A504B6A6D003A2000
+ - Name: .hip_fatbin
+ Type: SHT_PROGBITS
+ Flags: [ SHF_ALLOC ]
+ Address: 0x201000
+ AddressAlign: 0x1000
+ Content: 5F5F434C414E475F4F46464C4F41445F42554E444C455F5F0200000000000000001000000000000000000000000000001B00000000000000686F73742D7838365F36342D756E6B6E6F776E2D6C696E75782D2D0010000000000000D00F0000000000001F0000000000000068697076342D616D6467636E2D616D642D616D646873612D2D676678393038000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000007F454C460201014003000000000000000300E0000100000000000000000000004000000000000000100C0000000000003005000040003800090040000F000D000600000004000000400000000000000040000000000000004000000000000000F801000000000000F80100000000000008000000000000000100000004000000000000000000000000000000000000000000000000000000C008000000000000C008000000000000001000000000000001000000050000000009000000000000001900000000000000190000000000006C000000000000006C00000000000000001000000000000001000000060000007009000000000000702900000000000070290000000000007000000000000000900600000000000000100000000000000100000006000000E009000000000000E039000000000000E039000000000000000000000000000001000000000000000010000000000000020000000600000070090000000000007029000000000000702900000000000070000000000000007000000000000000080000000000000052E574640400000070090000000000007029000000000000702900000000000070000000000000009006000000000000010000000000000051E57464060000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000400000004000000380200000000000038020000000000003802000000000000340500000000000034050000000000000400000000000000070000001D05000020000000414D44475055000083AE616D646873612E6B65726E656C7391DE0012AB2E616770725F636F756E7400A52E61726773DC001085AE2E616464726573735F7370616365A6676C6F62616CA52E6E616D65AA415F642E636F65726365A72E6F666673657400A52E73697A6508AB2E76616C75655F6B696E64AD676C6F62616C5F62756666657285AE2E616464726573735F7370616365A6676C6F62616CA52E6E616D65AA425F642E636F65726365A72E6F666673657408A52E73697A6508AB2E76616C75655F6B696E64AD676C6F62616C5F62756666657284A52E6E616D65A14EA72E6F666673657410A52E73697A6508AB2E76616C75655F6B696E64A862795F76616C756583A72E6F666673657418A52E73697A6504AB2E76616C75655F6B696E64B468696464656E5F626C6F636B5F636F756E745F7883A72E6F66667365741CA52E73697A6504AB2E76616C75655F6B696E64B468696464656E5F626C6F636B5F636F756E745F7983A72E6F666673657420A52E73697A6504AB2E76616C75655F6B696E64B468696464656E5F626C6F636B5F636F756E745F7A83A72E6F666673657424A52E73697A6502AB2E76616C75655F6B696E64B368696464656E5F67726F75705F73697A655F7883A72E6F666673657426A52E73697A6502AB2E76616C75655F6B696E64B368696464656E5F67726F75705F73697A655F7983A72E6F666673657428A52E73697A6502AB2E76616C75655F6B696E64B368696464656E5F67726F75705F73697A655F7A83A72E6F66667365742AA52E73697A6502AB2E76616C75655F6B696E64B268696464656E5F72656D61696E6465725F7883A72E6F66667365742CA52E73697A6502AB2E76616C75655F6B696E64B268696464656E5F72656D61696E6465725F7983A72E6F66667365742EA52E73697A6502AB2E76616C75655F6B696E64B268696464656E5F72656D61696E6465725F7A83A72E6F666673657440A52E73697A6508AB2E76616C75655F6B696E64B668696464656E5F676C6F62616C5F6F66667365745F7883A72E6F666673657448A52E73697A6508AB2E76616C75655F6B696E64B668696464656E5F676C6F62616C5F6F66667365745F7983A72E6F666673657450A52E73697A6508AB2E76616C75655F6B696E64B668696464656E5F676C6F62616C5F6F66667365745F7A83A72E6F666673657458A52E73697A6502AB2E76616C75655F6B696E64B068696464656E5F677269645F64696D73B92E67726F75705F7365676D656E745F66697865645F73697A6500B62E6B65726E6172675F7365676D656E745F616C69676E08B52E6B65726E6172675F7365676D656E745F73697A65CD0118A92E6C616E6775616765A84F70656E434C2043B12E6C616E67756167655F76657273696F6E920200B82E6D61785F666C61745F776F726B67726F75705F73697A65CD0400A52E6E616D65B25F5A3973696D706C65416464506A504B6A6DBB2E707269766174655F7365676D656E745F66697865645F73697A6500AB2E736770725F636F756E740CB12E736770725F7370696C6C5F636F756E7400A72E73796D626F6CB55F5A3973696D706C65416464506A504B6A6D2E6B64B82E756E69666F726D5F776F726B5F67726F75705F73697A6501B32E757365735F64796E616D69635F737461636BC2AB2E766770725F636F756E7404B12E766770725F7370696C6C5F636F756E7400AF2E7761766566726F6E745F73697A6540AD616D646873612E746172676574B9616D6467636E2D616D642D616D646873612D2D676678393038AE616D646873612E76657273696F6E92010200000000000000000000000000000000000000000000000000000000000000010000001203070000190000000000006C000000000000001400000011030600800800000000000040000000000000002A00000011000A00E03900000000000001000000000000000100000001000000010000001A000000000008400000D20001000000360A4A7A5238A4D3F113F4DD04000000040000000200000001000000000000000300000000000000000000000000000000000000005F5A3973696D706C65416464506A504B6A6D005F5A3973696D706C65416464506A504B6A6D2E6B64005F5F6869705F637569645F623730363264386333326134613933330000000000000000000000000000000000000000000000000000000000000000000000180100000000000080100000000000000000000000000000000000000000000000000000000000004000AF008C000000090000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000C20102C02400000002000AC0000000008002027E7FC08CBF07FF0486FFFF0000060406920600006800008FD2820002000302067E0200043203030638008050DC02007F020102067E0000003203030238008050DC00007F03700F8CBF03050468008070DC00027F00000081BF00000000060000000000000070070000000000000B000000000000001800000000000000050000000000000020080000000000000A000000000000004600000000000000F5FEFF6F00000000D0070000000000000400000000000000F807000000000000000000000000000000000000000000004C696E6B65723A20414D44204C4C442031392E302E3000414D4420636C616E672076657273696F6E2031392E302E306769742028202032343231322063393630313665636534313337356462646438663037356266333762643666633333323230376233290000414D4420636C616E672076657273696F6E2031382E302E3067697420287373683A2F2F6765727269746769742F6C696768746E696E672F65632F6C6C766D2D70726F6A65637420616D642D6D61696E6C696E652D6F70656E20323431373620663935303039613166393032313232343865313036333964653837653635636163616338643961372900000000000000000000000000000000000000000000000000460000000002080070290000000000000000000000000000010000001203070000190000000000006C000000000000001400000011030600800800000000000040000000000000002A00000011000A00E0390000000000000100000000000000002E6E6F7465002E64796E73796D002E676E752E68617368002E68617368002E64796E737472002E726F64617461002E74657874002E64796E616D6963002E72656C726F5F70616464696E67002E627373002E636F6D6D656E74002E73796D746162002E7368737472746162002E73747274616200005F5A3973696D706C65416464506A504B6A6D005F5A3973696D706C65416464506A504B6A6D2E6B64005F5F6869705F637569645F62373036326438633332613461393333005F44594E414D494300000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000001000000070000000200000000000000380200000000000038020000000000003405000000000000000000000000000004000000000000000000000000000000070000000B00000002000000000000007007000000000000700700000000000060000000000000000500000001000000080000000000000018000000000000000F000000F6FFFF6F0200000000000000D007000000000000D007000000000000280000000000000002000000000000000800000000000000000000000000000019000000050000000200000000000000F807000000000000F80700000000000028000000000000000200000000000000040000000000000004000000000000001F000000030000000200000000000000200800000000000020080000000000004600000000000000000000000000000001000000000000000000000000000000270000000100000002000000000000008008000000000000800800000000000040000000000000000000000000000000400000000000000000000000000000002F000000010000000600000000000000001900000000000000090000000000006C00000000000000000000000000000000010000000000000000000000000000350000000600000003000000000000007029000000000000700900000000000070000000000000000500000000000000080000000000000010000000000000003E000000080000000300000000000000E029000000000000E00900000000000020060000000000000000000000000000010000000000000000000000000000004D000000080000000300000000000000E039000000000000E0090000000000000100000000000000000000000000000001000000000000000000000000000000520000000100000030000000000000000000000000000000E009000000000000F0000000000000000000000000000000010000000000000001000000000000005B0000000200000000000000000000000000000000000000D00A00000000000078000000000000000E0000000200000008000000000000001800000000000000630000000300000000000000000000000000000000000000480B00000000000075000000000000000000000000000000010000000000000000000000000000006D0000000300000000000000000000000000000000000000BD0B0000000000004F00000000000000000000000000000001000000000000000000000000000000
+ - Name: .hipFatBinSegment
+ Type: SHT_PROGBITS
+ Flags: [ SHF_ALLOC ]
+ Address: 0x202FD0
+ AddressAlign: 0x8
+ Content: '465049480100000000102000000000000000000000000000'
+ - Name: .eh_frame_hdr
+ Type: SHT_PROGBITS
+ Flags: [ SHF_ALLOC ]
+ Address: 0x202FE8
+ AddressAlign: 0x4
+ Content: 011B033B5C0000000A000000C811000078000000F81100009000000008120000A400000048120000C400000098120000E400000018130000040100004818000070010000A818000088010000D818000058010000F8180000A4010000
+ - Name: .eh_frame
+ Type: SHT_PROGBITS
+ Flags: [ SHF_ALLOC ]
+ Address: 0x203048
+ AddressAlign: 0x8
+ Content: 1400000000000000017A5200017810011B0C070890010000140000001C0000004811000026000000004407100000000010000000340000006011000005000000000000001C000000480000005C1100003900000000410E108602430D06740C07080000001C000000680000007C1100004C00000000410E108602430D0602470C070800001C00000088000000AC1100007900000000440E7002670E78440E8001490E080050000000A80000000C1200002C05000000420E10420E18420E20420E28410E30470ED00C83068C058D048E038F02037F030ED80C440EE00C490ED00C0381010E30410E28420E20420E18420E10420E08410ED00C14000000FC000000781700002000000000410E105A0E08001400000014010000D01600005B00000000440E3002520E08180000002C0100001817000022000000004D0E10540E08000000000014000000480100004C17000012000000000000000000000000000000
+ - Name: .text
+ Type: SHT_PROGBITS
+ Flags: [ SHF_ALLOC, SHF_EXECINSTR ]
+ Address: 0x2041B0
+ AddressAlign: 0x10
+ Content: F30F1EFA31ED4989D15E4889E24883E4F050544531C031C9488D3D31010000FF15FB1A0000F4662E0F1F840000000000F30F1EFAC3CCCCCCCCCCCCCCCCCCCCCC554889E5F605052C0000017402EB28C605FA2B000001488B05D31A00004885C07415EB00488D3D2DEEFFFF488D35E62B0000E8090700005DC30F1F8000000000554889E5F6050D2C0000017402EB3BC605022C000001488B059B1A00004885C0740EEB00488B3DA52A0000E8E0060000488B05891A00004885C0740EEB00488D3DD3EDFFFFE8D60600005DC3CCCCCCCC4883EC6848897C244848897424404889542438488D4424484889442450488D4424404889442458488D4424384889442460488D7C2428488D742418488D542410488D4C2408E896060000488B7424288B542430488B4C2418448B4424204C8D4C2450BF900D2000FF742408FF742418E87C0600004883C478C30F1F80000000004157415641554154534881EC20060000C744240400000000488D7C2404E85E060000BF806E2000BECA0D2000BA06000000E85A0600008B742404BF806E2000E85C0600004889C3BEA20D2000BA0D0000004889C7E837060000488B03488B40E84C8BB403F00000004D85F60F84B604000041807E38007407410FB64643EB164C89F7E829060000498B064C89F7BE0A000000FF50300FBEF04889DFE8200600004889C7E828060000837C2404000F8E2401000031DB4C8D742460EB3B0F1F40004C89FFE8E8050000498B074C89FFBE0A000000FF50300FBEF0BF806E2000E8DD0500004889C7E8E5050000FFC33B5C24040F8DE0000000BF806E2000BE980D2000BA07000000E885050000BF806E200089DEE889050000488B08488B49E84C8BBC08F00000004D85FF0F84F803000041807F3800740A410FB64F43EB210F1F004C89FF4989C4E865050000498B074C89FFBE0A000000FF503089C14C89E00FBEF14889C7E8570500004889C7E85F0500004C89F789DEE865050000BF806E2000BED10D2000BA07000000E8010500004C89F7E859050000BF806E20004C89F64889C2E8E9040000488B05D2290000488B40E84C8BB8706F20004D85FF0F846503000041807F38000F84F3FEFFFF410FB64743E9FFFEFFFF31FFE822050000BFA00F0000E8280500004889C3BFA00F0000E81B0500004989C6660F6F0550C8FFFFB80C000000660F6F0D53C8FFFF660F6F152BC8FFFF660F6F1D53C8FFFF660F6F255BC8FFFF66662E0F1F840000000000660F6FE8660FFEE9F30F7F4483D0F30F7F6C83E0660F6FF0660FFEF0660FFEEDF3410F7F7486D0F3410F7F6C86E0483DEC030000743A660F6FE8660FFEEA660F6FF0660FFEF3F30F7F6C83F0F30F7F3483660FFEED660FFEF6F3410F7F6C86F0F3410F7F3486660FFEC44883C010EB90488D7C2408BEA00F0000E871040000488D7C2410BEA00F0000E862040000488B7C2408BAA00F00004889DEB901000000E85B040000488B7C2410BAA00F00004C89F6B901000000E84404000048BA0001000001000000488DBA00010000BE01000000B9010000004531C04531C9E82E04000085C00F8582000000488B442408488B4C2410488944245848894C245048C7442448E8030000488D4424584889442460488D4424504889442468488D4424484889442470488D7C2438488D742428488D542420488D4C2418E8F2020000488B7424388B542440488B4C2428448B4424304C8D4C2460BF900D2000FF742418FF742428E8D80200004883C410488B742408BAA00F00004889DFB902000000E87D0300004531FF4531E40F1F8000000000428B34A3BF806E2000E882030000BEA00D2000BA010000004889C7E8B0020000463B3CA3751549FFC44983C7034981FCE803000075CAE9B9000000BF806E2000BEB00D2000BA14000000E881020000BF806E20004C89E6E8340300004989C5BEEC0D2000BA020000004889C7E85F020000428B34A34C89EFE8130300004989C4BEC50D2000BA040000004889C7E83E0200004C89E74C89FEE8F3020000488B08488B49E84C8BBC08F00000004D85FF0F84B200000041807F38007407410FB64F43EB1E4C89FF4989C4E822020000498B074C89FFBE0A000000FF503089C14C89E00FBEF14889C7E8140200004889C7E81C020000488B05C5260000488B40E84C8BB8706F20004D85FF745C41807F38007407410FB64743EB164C89FFE8CF010000498B074C89FFBE0A000000FF50300FBEF0BF806E2000E8C40100004889C7E8CC0100004889DFE8540200004C89F7E84C02000031C04881C4200600005B415C415D415E415FC3E8440200000F1F40004883EC28488B3D1D2600004885FF7514BFD02F2000E8660200004889C7488905042600000F57C00F114424100F110424BE900D2000BAD90D2000B9D90D200041B8FFFFFFFF4531C9E843020000BF904820004883C428E9550000000F1F440000488B3DC12500004885FF741550E82E02000048C705AB250000000000004883C408C3CCCCCCCCCCCCCCCCCCCCCCCCCCCC50BF506E2000E8B5010000BF904A2000BE506E2000BA006D200058E9C0010000F30F1EFA488B151524000031F6E9AE010000CCCC
+ - Name: .init
+ Type: SHT_PROGBITS
+ Flags: [ SHF_ALLOC, SHF_EXECINSTR ]
+ Address: 0x2048F4
+ AddressAlign: 0x4
+ Content: F30F1EFA4883EC08488B05D51300004885C07402FFD04883C408C3
+ - Name: .fini
+ Type: SHT_PROGBITS
+ Flags: [ SHF_ALLOC, SHF_EXECINSTR ]
+ Address: 0x204910
+ AddressAlign: 0x4
+ Content: F30F1EFA4883EC084883C408C3
+ - Name: .plt
+ Type: SHT_PROGBITS
+ Flags: [ SHF_ALLOC, SHF_EXECINSTR ]
+ Address: 0x204920
+ AddressAlign: 0x10
+ Content: FF35EA230000FF25EC2300000F1F4000FF25EA2300006800000000E9E0FFFFFFFF25E22300006801000000E9D0FFFFFFFF25DA2300006802000000E9C0FFFFFFFF25D22300006803000000E9B0FFFFFFFF25CA2300006804000000E9A0FFFFFFFF25C22300006805000000E990FFFFFFFF25BA2300006806000000E980FFFFFFFF25B22300006807000000E970FFFFFFFF25AA2300006808000000E960FFFFFFFF25A22300006809000000E950FFFFFFFF259A230000680A000000E940FFFFFFFF2592230000680B000000E930FFFFFFFF258A230000680C000000E920FFFFFFFF2582230000680D000000E910FFFFFFFF257A230000680E000000E900FFFFFFFF2572230000680F000000E9F0FEFFFFFF256A2300006810000000E9E0FEFFFFFF25622300006811000000E9D0FEFFFFFF255A2300006812000000E9C0FEFFFFFF25522300006813000000E9B0FEFFFFFF254A2300006814000000E9A0FEFFFFFF25422300006815000000E990FEFFFFFF253A2300006816000000E980FEFFFFFF25322300006817000000E970FEFFFFFF252A2300006818000000E960FEFFFFFF25222300006819000000E950FEFFFFFF251A230000681A000000E940FEFFFF
+ - Name: .init_array
+ Type: SHT_INIT_ARRAY
+ Flags: [ SHF_WRITE, SHF_ALLOC ]
+ Address: 0x205AE0
+ AddressAlign: 0x8
+ Content: F041200000000000C0482000000000003048200000000000
+ - Name: .fini_array
+ Type: SHT_FINI_ARRAY
+ Flags: [ SHF_WRITE, SHF_ALLOC ]
+ Address: 0x205AF8
+ AddressAlign: 0x8
+ Content: '3042200000000000'
+ - Name: .dynamic
+ Type: SHT_DYNAMIC
+ Flags: [ SHF_WRITE, SHF_ALLOC ]
+ Address: 0x205B00
+ Link: .dynstr
+ AddressAlign: 0x8
+ Entries:
+ - Tag: DT_RUNPATH
+ Value: 0x2E7
+ - Tag: DT_NEEDED
+ Value: 0x257
+ - Tag: DT_NEEDED
+ Value: 0x278
+ - Tag: DT_NEEDED
+ Value: 0x316
+ - Tag: DT_NEEDED
+ Value: 0x2B0
+ - Tag: DT_NEEDED
+ Value: 0x2C6
+ - Tag: DT_DEBUG
+ Value: 0x0
+ - Tag: DT_RELA
+ Value: 0x200A18
+ - Tag: DT_RELASZ
+ Value: 0x90
+ - Tag: DT_RELAENT
+ Value: 0x18
+ - Tag: DT_JMPREL
+ Value: 0x200AA8
+ - Tag: DT_PLTRELSZ
+ Value: 0x288
+ - Tag: DT_PLTGOT
+ Value: 0x206D08
+ - Tag: DT_PLTREL
+ Value: 0x7
+ - Tag: DT_SYMTAB
+ Value: 0x2002E8
+ - Tag: DT_SYMENT
+ Value: 0x18
+ - Tag: DT_STRTAB
+ Value: 0x2006F4
+ - Tag: DT_STRSZ
+ Value: 0x320
+ - Tag: DT_GNU_HASH
+ Value: 0x2006D0
+ - Tag: DT_INIT_ARRAY
+ Value: 0x205AE0
+ - Tag: DT_INIT_ARRAYSZ
+ Value: 0x18
+ - Tag: DT_FINI_ARRAY
+ Value: 0x205AF8
+ - Tag: DT_FINI_ARRAYSZ
+ Value: 0x8
+ - Tag: DT_INIT
+ Value: 0x2048F4
+ - Tag: DT_FINI
+ Value: 0x204910
+ - Tag: DT_VERSYM
+ Value: 0x2005D0
+ - Tag: DT_VERNEED
+ Value: 0x200610
+ - Tag: DT_VERNEEDNUM
+ Value: 0x4
+ - Tag: DT_NULL
+ Value: 0x0
+ - Name: .got
+ Type: SHT_PROGBITS
+ Flags: [ SHF_WRITE, SHF_ALLOC ]
+ Address: 0x205CD0
+ AddressAlign: 0x8
+ Content: '00000000000000000000000000000000000000000000000000000000000000000000000000000000'
+ - Name: .relro_padding
+ Type: SHT_NOBITS
+ Flags: [ SHF_WRITE, SHF_ALLOC ]
+ Address: 0x205CF8
+ AddressAlign: 0x1
+ Size: 0x308
+ - Name: .data
+ Type: SHT_PROGBITS
+ Flags: [ SHF_WRITE, SHF_ALLOC ]
+ Address: 0x206CF8
+ AddressAlign: 0x8
+ Content: 0000000000000000006D200000000000
+ - Name: .got.plt
+ Type: SHT_PROGBITS
+ Flags: [ SHF_WRITE, SHF_ALLOC ]
+ Address: 0x206D08
+ AddressAlign: 0x8
+ Content: 005B200000000000000000000000000000000000000000003649200000000000464920000000000056492000000000006649200000000000764920000000000086492000000000009649200000000000A649200000000000B649200000000000C649200000000000D649200000000000E649200000000000F649200000000000064A200000000000164A200000000000264A200000000000364A200000000000464A200000000000564A200000000000664A200000000000764A200000000000864A200000000000964A200000000000A64A200000000000B64A200000000000C64A200000000000D64A200000000000
+ - Name: .bss
+ Type: SHT_NOBITS
+ Flags: [ SHF_WRITE, SHF_ALLOC ]
+ Address: 0x206E00
+ AddressAlign: 0x40
+ Offset: 0x3DF8
+ Size: 0x190
+ - Name: .comment
+ Type: SHT_PROGBITS
+ Flags: [ SHF_MERGE, SHF_STRINGS ]
+ AddressAlign: 0x1
+ EntSize: 0x1
+ Content: 4C696E6B65723A20414D44204C4C442031392E302E3000414D4420636C616E672076657273696F6E2031392E302E306769742028202032343231322063393630313665636534313337356462646438663037356266333762643666633333323230376233290000414D4420636C616E672076657273696F6E2031392E302E3067697420282020323431393320373139633463633762336363396237353535333365363639656439316435373935346437373336352900
+Symbols:
+ - Name: __abi_tag
+ Type: STT_OBJECT
+ Section: .note.ABI-tag
+ Value: 0x2002C4
+ Size: 0x20
+ - Name: _dl_relocate_static_pie
+ Type: STT_FUNC
+ Section: .text
+ Value: 0x2041E0
+ Size: 0x5
+ Other: [ STV_HIDDEN ]
+ - Name: crtbegin.c
+ Type: STT_FILE
+ Index: SHN_ABS
+ - Name: __do_init
+ Type: STT_FUNC
+ Section: .text
+ Value: 0x2041F0
+ Size: 0x39
+ - Name: __do_init.__initialized
+ Type: STT_OBJECT
+ Section: .bss
+ Value: 0x206E00
+ Size: 0x1
+ - Name: __EH_FRAME_LIST__
+ Type: STT_OBJECT
+ Section: .eh_frame
+ Value: 0x203048
+ - Name: __do_init.__object
+ Type: STT_OBJECT
+ Section: .bss
+ Value: 0x206E08
+ Size: 0x40
+ - Name: __do_fini
+ Type: STT_FUNC
+ Section: .text
+ Value: 0x204230
+ Size: 0x4C
+ - Name: __do_fini.__finalized
+ Type: STT_OBJECT
+ Section: .bss
+ Value: 0x206E48
+ Size: 0x1
+ - Name: __init
+ Type: STT_OBJECT
+ Section: .init_array
+ Value: 0x205AE0
+ Size: 0x8
+ - Name: __fini
+ Type: STT_OBJECT
+ Section: .fini_array
+ Value: 0x205AF8
+ Size: 0x8
+ - Name: __dso_handle
+ Type: STT_OBJECT
+ Section: .data
+ Value: 0x206D00
+ Size: 0x8
+ Other: [ STV_HIDDEN ]
+ - Name: simpleAdd.cpp
+ Type: STT_FILE
+ Index: SHN_ABS
+ - Name: _GLOBAL__sub_I_simpleAdd.cpp
+ Type: STT_FUNC
+ Section: .text
+ Value: 0x2048C0
+ Size: 0x20
+ - Name: _ZStL8__ioinit
+ Type: STT_OBJECT
+ Section: .bss
+ Value: 0x206E50
+ Size: 0x1
+ - Name: __hip_module_ctor
+ Type: STT_FUNC
+ Section: .text
+ Value: 0x204830
+ Size: 0x5B
+ - Name: __hip_gpubin_handle
+ Type: STT_OBJECT
+ Section: .bss
+ Value: 0x206E58
+ Size: 0x8
+ - Name: __hip_fatbin_wrapper
+ Type: STT_OBJECT
+ Section: .hipFatBinSegment
+ Value: 0x202FD0
+ Size: 0x18
+ - Name: __hip_module_dtor
+ Type: STT_FUNC
+ Section: .text
+ Value: 0x204890
+ Size: 0x22
+ - Name: crtend.c
+ Type: STT_FILE
+ Index: SHN_ABS
+ - Name: __EH_FRAME_LIST_END__
+ Type: STT_OBJECT
+ Section: .eh_frame
+ Value: 0x203048
+ Size: 0x4
+ Other: [ STV_HIDDEN ]
+ - Name: _GLOBAL_OFFSET_TABLE_
+ Section: .got.plt
+ Value: 0x206D08
+ Other: [ STV_HIDDEN ]
+ - Name: _DYNAMIC
+ Section: .dynamic
+ Value: 0x205B00
+ Other: [ STV_HIDDEN ]
+ - Name: _init
+ Type: STT_FUNC
+ Section: .init
+ Value: 0x2048F4
+ Other: [ STV_HIDDEN ]
+ - Name: _fini
+ Type: STT_FUNC
+ Section: .fini
+ Value: 0x204910
+ Other: [ STV_HIDDEN ]
+ - Name: atexit
+ Type: STT_FUNC
+ Section: .text
+ Value: 0x2048E0
+ Size: 0x12
+ Other: [ STV_HIDDEN ]
+ - Name: _start
+ Type: STT_FUNC
+ Section: .text
+ Binding: STB_GLOBAL
+ Value: 0x2041B0
+ Size: 0x26
+ - Name: main
+ Type: STT_FUNC
+ Section: .text
+ Binding: STB_GLOBAL
+ Value: 0x204300
+ Size: 0x52C
+ - Name: data_start
+ Section: .data
+ Binding: STB_WEAK
+ Value: 0x206CF8
+ - Name: _IO_stdin_used
+ Type: STT_OBJECT
+ Section: .rodata
+ Binding: STB_GLOBAL
+ Value: 0x200D30
+ Size: 0x4
+ - Name: __libc_start_main
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: __data_start
+ Section: .data
+ Binding: STB_GLOBAL
+ Value: 0x206CF8
+ - Name: __gmon_start__
+ Binding: STB_WEAK
+ - Name: __register_frame_info
+ Type: STT_FUNC
+ Binding: STB_WEAK
+ - Name: __cxa_finalize
+ Type: STT_FUNC
+ Binding: STB_WEAK
+ - Name: __deregister_frame_info
+ Type: STT_FUNC
+ Binding: STB_WEAK
+ - Name: _Z24__device_stub__simpleAddPjPKjm
+ Type: STT_FUNC
+ Section: .text
+ Binding: STB_GLOBAL
+ Value: 0x204280
+ Size: 0x79
+ - Name: __hipPopCallConfiguration
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: _Z9simpleAddPjPKjm
+ Type: STT_OBJECT
+ Section: .rodata
+ Binding: STB_GLOBAL
+ Value: 0x200D90
+ Size: 0x8
+ - Name: hipLaunchKernel
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: hipGetDeviceCount
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: _ZSt4cout
+ Type: STT_OBJECT
+ Section: .bss
+ Binding: STB_GLOBAL
+ Value: 0x206E80
+ Size: 0x110
+ - Name: _ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: _ZNSolsEi
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: _ZNKSt5ctypeIcE13_M_widen_initEv
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: _ZNSo3putEc
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: _ZNSo5flushEv
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: hipGetDevicePropertiesR0600
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: strlen
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: hipSetDevice
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: malloc
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: hipMalloc
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: hipMemcpy
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: __hipPushCallConfiguration
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: _ZNSo9_M_insertImEERSoT_
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: free
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: _ZSt16__throw_bad_castv
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: _ZNSt8ios_base4InitC1Ev
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: _ZNSt8ios_base4InitD1Ev
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ Value: 0x204A90
+ - Name: __cxa_atexit
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: __hipRegisterFatBinary
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: __hipRegisterFunction
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: __hipUnregisterFatBinary
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: __hip_cuid_b7062d8c32a4a933
+ Type: STT_OBJECT
+ Section: .bss
+ Binding: STB_GLOBAL
+ Value: 0x206E60
+ Size: 0x1
+DynamicSymbols:
+ - Name: __libc_start_main
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: __gmon_start__
+ Binding: STB_WEAK
+ - Name: __register_frame_info
+ Type: STT_FUNC
+ Binding: STB_WEAK
+ - Name: __cxa_finalize
+ Type: STT_FUNC
+ Binding: STB_WEAK
+ - Name: __deregister_frame_info
+ Type: STT_FUNC
+ Binding: STB_WEAK
+ - Name: __hipPopCallConfiguration
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: hipLaunchKernel
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: hipGetDeviceCount
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: _ZSt16__ostream_insertIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_PKS3_l
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: _ZNSolsEi
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: _ZNKSt5ctypeIcE13_M_widen_initEv
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: _ZNSo3putEc
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: _ZNSo5flushEv
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: hipGetDevicePropertiesR0600
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: strlen
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: hipSetDevice
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: malloc
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: hipMalloc
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: hipMemcpy
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: __hipPushCallConfiguration
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: _ZNSo9_M_insertImEERSoT_
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: free
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: _ZSt16__throw_bad_castv
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: _ZNSt8ios_base4InitC1Ev
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: __cxa_atexit
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: __hipRegisterFatBinary
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: __hipRegisterFunction
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: __hipUnregisterFatBinary
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ - Name: _ZSt4cout
+ Type: STT_OBJECT
+ Section: .bss
+ Binding: STB_GLOBAL
+ Value: 0x206E80
+ Size: 0x110
+ - Name: _ZNSt8ios_base4InitD1Ev
+ Type: STT_FUNC
+ Binding: STB_GLOBAL
+ Value: 0x204A90
+...
diff --git a/llvm/tools/llvm-objdump/ObjdumpOpts.td b/llvm/tools/llvm-objdump/ObjdumpOpts.td
index c3764c6e97534a..be0553d0f75d74 100644
--- a/llvm/tools/llvm-objdump/ObjdumpOpts.td
+++ b/llvm/tools/llvm-objdump/ObjdumpOpts.td
@@ -108,6 +108,9 @@ def fault_map_section : Flag<["--"], "fault-map-section">,
def offloading : Flag<["--"], "offloading">,
HelpText<"Display the content of the offloading section">;
+def offload_fatbin : Flag<["--"], "offload-fatbin">,
+ HelpText<"Display the content of the offload FatBin section">;
+
def file_headers : Flag<["--"], "file-headers">,
HelpText<"Display the contents of the overall file header">;
def : Flag<["-"], "f">, Alias<file_headers>,
diff --git a/llvm/tools/llvm-objdump/OffloadDump.cpp b/llvm/tools/llvm-objdump/OffloadDump.cpp
index 4ac6b99e79bbbb..516f5d02785441 100644
--- a/llvm/tools/llvm-objdump/OffloadDump.cpp
+++ b/llvm/tools/llvm-objdump/OffloadDump.cpp
@@ -14,12 +14,15 @@
#include "OffloadDump.h"
#include "llvm-objdump.h"
#include "llvm/Object/ELFObjectFile.h"
+#include "llvm/Object/OffloadBinary.h"
#include "llvm/Support/Alignment.h"
using namespace llvm;
using namespace llvm::object;
using namespace llvm::objdump;
+void disassembleObject(llvm::object::ObjectFile *, bool InlineRelocs);
+
/// Get the printable name of the image kind.
static StringRef getImageName(const OffloadBinary &OB) {
switch (OB.getImageKind()) {
@@ -66,6 +69,55 @@ void llvm::dumpOffloadBinary(const ObjectFile &O) {
printBinary(*Binaries[I].getBinary(), I);
}
+// Given an Object file, collect all Bundles of FatBin Binaries
+// and dump them into Code Object files
+// if -d is specified, disassemble the Code Object Files
+// if -arch=-name is specified, only dump the Entries that match the target arch
+void llvm::dumpOffloadFatBinary(const ObjectFile &O, std::string ArchName,
+ bool Disassemble) {
+ assert((O.isELF() || O.isCOFF()) && "Invalid file type");
+ // Collect all Bundles and their Entries ....
+ SmallVector<llvm::object::OffloadFatBinBundle> FoundBundles;
+ SmallVector<BundleEntry> FoundEntries;
+
+ if (Error Err = llvm::object::extractFatBinaryFromObject(O, FoundBundles))
+ reportError(O.getFileName(), "while extracting offload FatBin bundles: " +
+ toString(std::move(Err)));
+
+ // Now filter based on if arch-name is specified
+ SmallVectorImpl<llvm::object::OffloadFatBinBundle>::iterator BundleIter =
+ FoundBundles.begin();
+ for (uint64_t bundle_num = 0; bundle_num < FoundBundles.size();
+ bundle_num++) {
+ if (!ArchName.empty())
+ FoundEntries = BundleIter->EntryIDContains(StringRef(ArchName));
+ else
+ FoundEntries = BundleIter->getEntries();
+
+ // now we have a list of Found Entries .... dump them
+ SmallVectorImpl</*OffloadFatBinBundle::*/ BundleEntry>::iterator FoundIter =
+ FoundEntries.begin();
+ for (int64_t entry_num = 0; entry_num < FoundEntries.size(); entry_num++) {
+ // create file name for this object file: <source-filename>:<Bundle
+ // Number>.<EntryID>
+ std::string str = BundleIter->getFileName().str() + ":" +
+ itostr(bundle_num) + "." + FoundIter->ID.str();
+ StringRef OutputFilename = StringRef(str);
+ if (Error Err = object::extractCodeObject(
+ O, FoundIter->Offset, FoundIter->Size, OutputFilename))
+ reportError(O.getFileName(),
+ "while extracting offload Bundle Entries: " +
+ toString(std::move(Err)));
+
+ // TODO: If -d was specified, disasseble the Code Object too
+
+ ++FoundIter;
+ } // end of for found_entries loop
+
+ ++BundleIter;
+ } // end of for Bundles loop
+}
+
/// Print the contents of an offload binary file \p OB. This may contain
/// multiple binaries stored in the same buffer.
void llvm::dumpOffloadSections(const OffloadBinary &OB) {
diff --git a/llvm/tools/llvm-objdump/OffloadDump.h b/llvm/tools/llvm-objdump/OffloadDump.h
index 75f188e9d50656..9ba08183558009 100644
--- a/llvm/tools/llvm-objdump/OffloadDump.h
+++ b/llvm/tools/llvm-objdump/OffloadDump.h
@@ -16,7 +16,8 @@ namespace llvm {
void dumpOffloadSections(const object::OffloadBinary &OB);
void dumpOffloadBinary(const object::ObjectFile &O);
-
+void dumpOffloadFatBinary(const object::ObjectFile &O, std::string ArchName,
+ bool Disassemble);
} // namespace llvm
#endif
diff --git a/llvm/tools/llvm-objdump/llvm-objdump.cpp b/llvm/tools/llvm-objdump/llvm-objdump.cpp
index 86ba9193dff2d1..2d9894ef5a14e9 100644
--- a/llvm/tools/llvm-objdump/llvm-objdump.cpp
+++ b/llvm/tools/llvm-objdump/llvm-objdump.cpp
@@ -324,6 +324,7 @@ std::vector<std::string> objdump::MAttrs;
bool objdump::ShowRawInsn;
bool objdump::LeadingAddr;
static bool Offloading;
+static bool OffloadFatBin;
static bool RawClangAST;
bool objdump::Relocations;
bool objdump::PrintImmHex;
@@ -2529,7 +2530,7 @@ disassembleObject(ObjectFile &Obj, const ObjectFile &DbgObj,
reportWarning("failed to disassemble missing symbol " + Sym, FileName);
}
-static void disassembleObject(ObjectFile *Obj, bool InlineRelocs) {
+void disassembleObject(llvm::object::ObjectFile *Obj, bool InlineRelocs) {
// If information useful for showing the disassembly is missing, try to find a
// more complete binary and disassemble that instead.
OwningBinary<Binary> FetchedBinary;
@@ -3315,6 +3316,8 @@ static void dumpObject(ObjectFile *O, const Archive *A = nullptr,
D.printDynamicRelocations();
if (SectionContents)
printSectionContents(O);
+ if (OffloadFatBin)
+ dumpOffloadFatBinary(*O, ArchName, Disassemble);
if (Disassemble)
disassembleObject(O, Relocations);
if (UnwindInfo)
@@ -3521,6 +3524,7 @@ static void parseObjdumpOptions(const llvm::opt::InputArgList &InputArgs) {
DynamicRelocations = InputArgs.hasArg(OBJDUMP_dynamic_reloc);
FaultMapSection = InputArgs.hasArg(OBJDUMP_fault_map_section);
Offloading = InputArgs.hasArg(OBJDUMP_offloading);
+ OffloadFatBin = InputArgs.hasArg(OBJDUMP_offload_fatbin);
FileHeaders = InputArgs.hasArg(OBJDUMP_file_headers);
SectionContents = InputArgs.hasArg(OBJDUMP_full_contents);
PrintLines = InputArgs.hasArg(OBJDUMP_line_numbers);
@@ -3732,6 +3736,7 @@ int llvm_objdump_main(int argc, char **argv, const llvm::ToolContext &) {
!DynamicRelocations && !FileHeaders && !PrivateHeaders && !RawClangAST &&
!Relocations && !SectionHeaders && !SectionContents && !SymbolTable &&
!DynamicSymbolTable && !UnwindInfo && !FaultMapSection && !Offloading &&
+ !OffloadFatBin &&
!(MachOOpt &&
(Bind || DataInCode || ChainedFixups || DyldInfo || DylibId ||
DylibsUsed || ExportsTrie || FirstPrivateHeader ||
More information about the llvm-commits
mailing list