[llvm] Extend llvm objdump fatbin (PR #114834)
David Salinas via llvm-commits
llvm-commits at lists.llvm.org
Mon Nov 4 09:23:59 PST 2024
https://github.com/david-salinas created https://github.com/llvm/llvm-project/pull/114834
Utilize the new extensions to the LLVM Offloading API to extend to llvm-objdump to handle dumping fatbin offload bundles generated by HIP. This extension to llvm-objdump adds the option --offload-fatbin. Specifying this option will take the input object/executable and extract all offload fatbin bundle entries into distinct code object files with names reflecting the source file name combined with the Bundle Entry ID. Users can also use the --arch-name option to filter offload fatbin bundle entries by their target triple.
>From 71c9c5cb43c750ce35136b183d28a8a138d6f6e5 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/2] 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 | 142 ++++++++
llvm/lib/Object/OffloadBinary.cpp | 441 +++++++++++++++++++++++
2 files changed, 583 insertions(+)
diff --git a/llvm/include/llvm/Object/OffloadBinary.h b/llvm/include/llvm/Object/OffloadBinary.h
index c02aec8d956ed6..c63ef4824bc7a4 100644
--- a/llvm/include/llvm/Object/OffloadBinary.h
+++ b/llvm/include/llvm/Object/OffloadBinary.h
@@ -17,10 +17,12 @@
#ifndef LLVM_OBJECT_OFFLOADBINARY_H
#define LLVM_OBJECT_OFFLOADBINARY_H
+#include "llvm/Support/Compression.h"
#include "llvm/ADT/MapVector.h"
#include "llvm/ADT/SmallString.h"
#include "llvm/ADT/StringRef.h"
#include "llvm/Object/Binary.h"
+#include "llvm/Object/ObjectFile.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,126 @@ 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..75085b96a960bf 100644
--- a/llvm/lib/Object/OffloadBinary.cpp
+++ b/llvm/lib/Object/OffloadBinary.cpp
@@ -17,6 +17,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"
@@ -25,12 +26,18 @@
#include "llvm/Support/Alignment.h"
#include "llvm/Support/FileOutputBuffer.h"
#include "llvm/Support/SourceMgr.h"
+#include "llvm/Support/Timer.h"
+#include "llvm/Support/BinaryStreamReader.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,42 @@ 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 +213,95 @@ 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 +431,94 @@ 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 +602,224 @@ 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 2f18112f00908a2d343359252449c2da8f65b060 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 2/2] 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 | 53 +-
llvm/lib/Object/ObjectFile.cpp | 1 -
llvm/lib/Object/OffloadBinary.cpp | 7 +-
.../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 | 42 +
llvm/tools/llvm-objdump/OffloadDump.h | 2 +-
llvm/tools/llvm-objdump/llvm-objdump.cpp | 10 +-
9 files changed, 1015 insertions(+), 27 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 c63ef4824bc7a4..984f40cc194fe0 100644
--- a/llvm/include/llvm/Object/OffloadBinary.h
+++ b/llvm/include/llvm/Object/OffloadBinary.h
@@ -210,15 +210,6 @@ 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;
@@ -226,39 +217,66 @@ class OffloadFatBinBundle {
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 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:
+
+
+ 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;
}
}
+ 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;
+ }
+
OffloadFatBinBundle(MemoryBufferRef Source, StringRef file) :
FileName(file) {
NumberOfEntries = 0;
- Entries = std::make_unique<SmallVector<BundleEntry>>();
+ Entries = SmallVector<BundleEntry>();
}
};
@@ -323,11 +341,8 @@ struct OffloadBundleURI {
/// 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.
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 75085b96a960bf..b2288e8a47be28 100644
--- a/llvm/lib/Object/OffloadBinary.cpp
+++ b/llvm/lib/Object/OffloadBinary.cpp
@@ -259,9 +259,9 @@ Error OffloadFatBinBundle::ReadEntries(StringRef Buffer, uint64_t SectionOffset)
}
// 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();
@@ -287,7 +287,7 @@ Expected<std::unique_ptr<OffloadFatBinBundle>> OffloadFatBinBundle::create(Memor
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) {
@@ -450,7 +450,6 @@ Error object::extractFatBinaryFromObject(const ObjectFile &Obj, SmallVectorImpl<
} 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..120cc26f4fcaac 100644
--- a/llvm/tools/llvm-objdump/OffloadDump.cpp
+++ b/llvm/tools/llvm-objdump/OffloadDump.cpp
@@ -15,11 +15,14 @@
#include "llvm-objdump.h"
#include "llvm/Object/ELFObjectFile.h"
#include "llvm/Support/Alignment.h"
+#include "llvm/Object/OffloadBinary.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,45 @@ 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..0f044b85ec4cb4 100644
--- a/llvm/tools/llvm-objdump/OffloadDump.h
+++ b/llvm/tools/llvm-objdump/OffloadDump.h
@@ -16,7 +16,7 @@ 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..49d863cb52a4f3 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);
@@ -3731,7 +3735,7 @@ int llvm_objdump_main(int argc, char **argv, const llvm::ToolContext &) {
if (!ArchiveHeaders && !Disassemble && DwarfDumpType == DIDT_Null &&
!DynamicRelocations && !FileHeaders && !PrivateHeaders && !RawClangAST &&
!Relocations && !SectionHeaders && !SectionContents && !SymbolTable &&
- !DynamicSymbolTable && !UnwindInfo && !FaultMapSection && !Offloading &&
+ !DynamicSymbolTable && !UnwindInfo && !FaultMapSection && !Offloading && !OffloadFatBin &&
!(MachOOpt &&
(Bind || DataInCode || ChainedFixups || DyldInfo || DylibId ||
DylibsUsed || ExportsTrie || FirstPrivateHeader ||
@@ -3750,3 +3754,5 @@ int llvm_objdump_main(int argc, char **argv, const llvm::ToolContext &) {
return EXIT_SUCCESS;
}
+
+
More information about the llvm-commits
mailing list