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

Author: dsalinas

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;
   }
 };
 

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())

llvm/lib/Object/OffloadBinary.cpp

@@ -18,7 +18,6 @@
 #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"
@@ -267,10 +266,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 +297,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 +465,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);
         }
       }
 

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 );
+}