Add support for code object URI to ROCr

Adds the following:
	- New factory method to create a code object reader from
          file with offset and size.
	- A pair of queries on a loaded code object to get the URI name/length.
	- A bump to the AMD vendor loader extension API and its associated table.

Change-Id: I17c83e9c2447d29a43c438459395365f786a3611


[ROCm/ROCR-Runtime commit: 9eb735ec24]

projects/rocr-runtime/runtime/hsa-runtime/CMakeLists.txt

@@ -138,6 +138,7 @@ set ( SRCS "core/util/lnx/os_linux.cpp"
             "core/runtime/amd_blit_sdma.cpp"
             "core/runtime/amd_cpu_agent.cpp"
             "core/runtime/amd_gpu_agent.cpp"
+            "core/runtime/amd_hsa_loader.cpp"
             "core/runtime/amd_aql_queue.cpp"
             "core/runtime/amd_loader_context.cpp"
             "core/runtime/hsa_ven_amd_loader.cpp"

projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_hsa_loader.hpp

@@ -53,6 +53,16 @@
 #include <mutex>
 #include <vector>
 
+#if defined(_WIN32) || defined(_WIN64)
+#include <io.h>
+#define __read__  _read
+#define __lseek__ _lseek
+#else
+#include <unistd.h>
+#define __read__  read
+#define __lseek__ lseek
+#endif  // _WIN32 || _WIN64
+
 /// @brief Major version of the AMD HSA Loader. Major versions are not backwards
 /// compatible.
 #define AMD_HSA_LOADER_VERSION_MAJOR 0
@@ -96,6 +106,60 @@ namespace amd {
 namespace hsa {
 namespace loader {
 
+/// @class CodeObjectReaderWrapper.
+/// @brief Code Object Reader Wrapper.
+struct CodeObjectReaderWrapper final {
+ private:
+  std::string GetUriFromFile(int Fd, size_t Offset, size_t Size) const;
+  std::string GetUriFromMemoryBasic(const void *Mem, size_t Size) const;
+  std::string GetUriFromMemory(const void *Mem, size_t Size) const;
+
+ public:
+  /// @returns Handle equivalent of @p object.
+  static hsa_code_object_reader_t Handle(
+      const CodeObjectReaderWrapper *object) {
+    hsa_code_object_reader_t handle = {reinterpret_cast<uint64_t>(object)};
+    return handle;
+  }
+
+  /// @returns Object equivalent of @p handle.
+  static CodeObjectReaderWrapper *Object(
+      const hsa_code_object_reader_t &handle) {
+    CodeObjectReaderWrapper *object =
+      reinterpret_cast<CodeObjectReaderWrapper*>(handle.handle);
+    return object;
+  }
+
+  /// @brief Default constructor.
+  CodeObjectReaderWrapper(
+      const void *_code_object_memory, size_t _code_object_size,
+      size_t _code_object_offset, hsa_file_t _code_object_file_descriptor)
+    : code_object_memory(_code_object_memory)
+    , code_object_size(_code_object_size)
+    , code_object_offset(_code_object_offset)
+    , code_object_file_descriptor(_code_object_file_descriptor) {}
+
+  /// @brief Default destructor.
+  ~CodeObjectReaderWrapper() {}
+
+  bool ComesFromFile() {
+    return code_object_file_descriptor != -1;
+  }
+
+  std::string GetUri() {
+    if (ComesFromFile()) {
+      return GetUriFromFile(code_object_file_descriptor, code_object_offset, code_object_size);
+    } else {
+      return GetUriFromMemory(code_object_memory, code_object_size);
+    }
+  }
+
+  const void *code_object_memory;
+  size_t code_object_size;
+  size_t code_object_offset;
+  hsa_file_t code_object_file_descriptor;
+};
+
 //===----------------------------------------------------------------------===//
 // Context.                                                                   //
 //===----------------------------------------------------------------------===//
@@ -216,6 +280,7 @@ public:
   virtual uint64_t getLoadBase() const = 0;
   virtual uint64_t getLoadSize() const = 0;
   virtual int64_t getDelta() const = 0;
+  virtual std::string getUri() const = 0;
 
 protected:
   LoadedCodeObject() {}
@@ -291,6 +356,7 @@ public:
     hsa_agent_t agent,
     hsa_code_object_t code_object,
     const char *options,
+    const std::string &uri,
     hsa_loaded_code_object_t *loaded_code_object = nullptr) = 0;
 
   virtual hsa_status_t LoadCodeObject(
@@ -298,6 +364,7 @@ public:
     hsa_code_object_t code_object,
     size_t code_object_size,
     const char *options,
+    const std::string &uri,
     hsa_loaded_code_object_t *loaded_code_object = nullptr) = 0;
 
   virtual hsa_status_t Freeze(const char *options) = 0;

projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_hsa_loader.cpp

@@ -0,0 +1,175 @@
+////////////////////////////////////////////////////////////////////////////////
+//
+// The University of Illinois/NCSA
+// Open Source License (NCSA)
+//
+// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved.
+//
+// Developed by:
+//
+//                 AMD Research and AMD HSA Software Development
+//
+//                 Advanced Micro Devices, Inc.
+//
+//                 www.amd.com
+//
+// Permission is hereby granted, free of charge, to any person obtaining a copy
+// of this software and associated documentation files (the "Software"), to
+// deal with the Software without restriction, including without limitation
+// the rights to use, copy, modify, merge, publish, distribute, sublicense,
+// and/or sell copies of the Software, and to permit persons to whom the
+// Software is furnished to do so, subject to the following conditions:
+//
+//  - Redistributions of source code must retain the above copyright notice,
+//    this list of conditions and the following disclaimers.
+//  - Redistributions in binary form must reproduce the above copyright
+//    notice, this list of conditions and the following disclaimers in
+//    the documentation and/or other materials provided with the distribution.
+//  - Neither the names of Advanced Micro Devices, Inc,
+//    nor the names of its contributors may be used to endorse or promote
+//    products derived from this Software without specific prior written
+//    permission.
+//
+// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
+// THE CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR
+// OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE,
+// ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER
+// DEALINGS WITH THE SOFTWARE.
+//
+////////////////////////////////////////////////////////////////////////////////
+
+#include "core/inc/amd_hsa_loader.hpp"
+
+#include <linux/limits.h>
+#include <unistd.h>
+
+#include <cstring>
+#include <fstream>
+#include <iomanip>
+#include <sstream>
+#include <string>
+
+namespace {
+
+std::string EncodePathname(const char *Pathname) {
+  std::ostringstream ss;
+  unsigned char c;
+
+  ss.fill('0');
+  ss << "file://";
+
+  while ((c = *Pathname++) != '\0') {
+    if (isalnum(c) || c == '/' || c == '-' ||
+        c == '_' || c == '.' || c == '~') {
+      ss << c;
+    } else {
+      ss << std::uppercase;
+      ss << '%' << std::hex << std::setw(2) << static_cast<int>(c);
+      ss << std::nouppercase;
+    }
+  }
+
+  return ss.str();
+}
+
+}  // namespace
+
+namespace amd {
+namespace hsa {
+namespace loader {
+
+std::string CodeObjectReaderWrapper::GetUriFromFile(
+    int Fd, size_t Offset, size_t Size) const {
+#if !defined(_WIN32) && !defined(_WIN64)
+  std::ostringstream ProcFdPath;
+  ProcFdPath << "/proc/self/fd/" << Fd;
+
+  char FdPath[PATH_MAX];
+  memset(FdPath, 0, PATH_MAX);
+
+  if (readlink(ProcFdPath.str().c_str(), FdPath, PATH_MAX) == -1) {
+    return std::string();
+  }
+
+  std::ostringstream UriStream;
+  UriStream << EncodePathname(FdPath);
+  if (Size) {
+    UriStream << "#offset=" << Offset;
+    UriStream << "&size=" << Size;
+  }
+  return UriStream.str();
+#else
+  return std::string();
+#endif  // !defined(_WIN32) && !defined(_WIN64)
+}
+
+std::string CodeObjectReaderWrapper::GetUriFromMemoryBasic(
+    const void *Mem, size_t Size) const {
+  pid_t PID = getpid();
+  std::ostringstream UriStream;
+  UriStream << "memory://" << PID
+            << "#offset=0x" << std::hex << (uint64_t)Mem << std::dec
+            << "&size=" << Size;
+  return UriStream.str();
+}
+
+std::string CodeObjectReaderWrapper::GetUriFromMemory(
+    const void *Mem, size_t Size) const {
+#if !defined(_WIN32) && !defined(_WIN64)
+  std::ostringstream ProcMapsPath;
+  ProcMapsPath << "/proc/self/maps";
+
+  std::ifstream ProcMapsFile;
+  ProcMapsFile.open(ProcMapsPath.str(), std::ifstream::in);
+  if (!ProcMapsFile.is_open() || !ProcMapsFile.good()) {
+    return GetUriFromMemoryBasic(Mem, Size);
+  }
+
+  std::string ProcMapsLine;
+  while (std::getline(ProcMapsFile, ProcMapsLine)) {
+    std::stringstream TokenStream(ProcMapsLine);
+
+    uint64_t LowAddress, HighAddress;
+    char Dash;
+    TokenStream >> std::hex >> LowAddress >> std::dec
+                >> Dash
+                >> std::hex >> HighAddress >> std::dec;
+    if (Dash != '-') {
+      continue;
+    }
+
+    uint64_t MyAddress = reinterpret_cast<uint64_t>(Mem);
+    if (!(MyAddress >= LowAddress && MyAddress <= HighAddress)) {
+      continue;
+    }
+
+    std::string Perms, Dev, Pathname;
+    uint64_t Offset, INode;
+    TokenStream >> Perms
+                >> std::hex >> Offset >> std::dec
+                >> Dev
+                >> INode
+                >> Pathname;
+
+    if (INode == 0 || Pathname.empty()) {
+      return GetUriFromMemoryBasic(Mem, Size);
+    }
+
+    std::ostringstream UriStream;
+    UriStream << EncodePathname(Pathname.c_str());
+    UriStream << "#offset=" << Offset;
+    if (Size) {
+      UriStream << "&size=" << Size;
+    }
+    return UriStream.str();
+  }
+
+#endif  // !defined(_WIN32) && !defined(_WIN64)
+  return GetUriFromMemoryBasic(Mem, Size);
+}
+
+}  // namespace loader
+}  // namespace hsa
+}  // namespace amd

projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa.cpp

@@ -47,16 +47,6 @@
 #include <string>
 #include <sys/types.h>
 
-#if defined(_WIN32) || defined(_WIN64)
-#include <io.h>
-#define __read__  _read
-#define __lseek__ _lseek
-#else
-#include <unistd.h>
-#define __read__  read
-#define __lseek__ lseek
-#endif // _WIN32 || _WIN64
-
 #include "core/inc/runtime.h"
 #include "core/inc/agent.h"
 #include "core/inc/host_queue.h"
@@ -2054,45 +2044,12 @@ hsa_status_t hsa_code_object_iterate_symbols(
 //===--- Executable -------------------------------------------------------===//
 
 using common::Signed;
+using loader::CodeObjectReaderWrapper;
 using loader::Executable;
 using loader::Loader;
 
 namespace {
 
-/// @class CodeObjectReaderWrapper.
-/// @brief Code Object Reader Wrapper.
-struct CodeObjectReaderWrapper final : Signed<0x266E71EDBC718D2C> {
-  /// @returns Handle equivalent of @p object.
-  static hsa_code_object_reader_t Handle(
-      const CodeObjectReaderWrapper *object) {
-    hsa_code_object_reader_t handle = {reinterpret_cast<uint64_t>(object)};
-    return handle;
-  }
-
-  /// @returns Object equivalent of @p handle.
-  static CodeObjectReaderWrapper *Object(
-      const hsa_code_object_reader_t &handle) {
-    CodeObjectReaderWrapper *object = common::ObjectAt<CodeObjectReaderWrapper>(
-        handle.handle);
-    return object;
-  }
-
-  /// @brief Default constructor.
-  CodeObjectReaderWrapper(
-      const void *_code_object_memory, size_t _code_object_size,
-      bool _comes_from_file)
-    : code_object_memory(_code_object_memory)
-    , code_object_size(_code_object_size)
-    , comes_from_file(_comes_from_file) {}
-
-  /// @brief Default destructor.
-  ~CodeObjectReaderWrapper() {}
-
-  const void *code_object_memory;
-  const size_t code_object_size;
-  const bool comes_from_file;
-};
-
 Loader *GetLoader() {
   return core::Runtime::runtime_singleton_->loader();
 }
@@ -2111,6 +2068,10 @@ hsa_status_t hsa_code_object_reader_create_from_file(
     return HSA_STATUS_ERROR_INVALID_FILE;
   }
 
+  if (file_size == 0) {
+    return HSA_STATUS_ERROR_INVALID_CODE_OBJECT;
+  }
+
   if (__lseek__(file, 0, SEEK_SET) == (off_t)-1) {
     return HSA_STATUS_ERROR_INVALID_FILE;
   }
@@ -2124,7 +2085,7 @@ hsa_status_t hsa_code_object_reader_create_from_file(
   }
 
   CodeObjectReaderWrapper *wrapper = new (std::nothrow) CodeObjectReaderWrapper(
-      code_object_memory, file_size, true);
+      code_object_memory, file_size, 0, file);
   if (!wrapper) {
     delete [] code_object_memory;
     return HSA_STATUS_ERROR_OUT_OF_RESOURCES;
@@ -2149,7 +2110,7 @@ hsa_status_t hsa_code_object_reader_create_from_memory(
   }
 
   CodeObjectReaderWrapper *wrapper = new (std::nothrow) CodeObjectReaderWrapper(
-      code_object, size, false);
+      code_object, size, 0, -1);
   CHECK_ALLOC(wrapper);
 
   *code_object_reader = CodeObjectReaderWrapper::Handle(wrapper);
@@ -2168,7 +2129,7 @@ hsa_status_t hsa_code_object_reader_destroy(
     return HSA_STATUS_ERROR_INVALID_CODE_OBJECT_READER;
   }
 
-  if (wrapper->comes_from_file) {
+  if (wrapper->ComesFromFile()) {
     delete [] (unsigned char*)wrapper->code_object_memory;
   }
   delete wrapper;
@@ -2261,7 +2222,7 @@ hsa_status_t hsa_executable_load_code_object(
     return HSA_STATUS_ERROR_INVALID_EXECUTABLE;
   }
 
-  return exec->LoadCodeObject(agent, code_object, options);
+  return exec->LoadCodeObject(agent, code_object, options, std::string());
   CATCH;
 }
 
@@ -2287,7 +2248,7 @@ hsa_status_t hsa_executable_load_program_code_object(
   hsa_code_object_t code_object =
       {reinterpret_cast<uint64_t>(wrapper->code_object_memory)};
   return exec->LoadCodeObject(
-      {0}, code_object, options, loaded_code_object);
+      {0}, code_object, options, wrapper->GetUri(), loaded_code_object);
   CATCH;
 }
 
@@ -2314,7 +2275,7 @@ hsa_status_t hsa_executable_load_agent_code_object(
   hsa_code_object_t code_object =
       {reinterpret_cast<uint64_t>(wrapper->code_object_memory)};
   return exec->LoadCodeObject(
-      agent, code_object, options, loaded_code_object);
+      agent, code_object, options, wrapper->GetUri(), loaded_code_object);
   CATCH;
 }
 

projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa_ven_amd_loader.cpp

@@ -48,6 +48,7 @@
 using namespace amd::hsa;
 using namespace core;
 
+using loader::CodeObjectReaderWrapper;
 using loader::Executable;
 using loader::LoadedCodeObject;
 
@@ -205,6 +206,14 @@ hsa_status_t hsa_ven_amd_loader_loaded_code_object_get_info(
       *((uint64_t*)value) = lcobj->getLoadSize();
       break;
     }
+    case HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_URI_LENGTH: {
+      *(reinterpret_cast<uint32_t*>(value)) = lcobj->getUri().size();
+      break;
+    }
+    case HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_URI: {
+      memcpy(value, lcobj->getUri().c_str(), lcobj->getUri().size());
+      break;
+    }
     default: {
       return HSA_STATUS_ERROR_INVALID_ARGUMENT;
     }
@@ -212,3 +221,45 @@ hsa_status_t hsa_ven_amd_loader_loaded_code_object_get_info(
 
   return HSA_STATUS_SUCCESS;
 }
+
+hsa_status_t
+hsa_ven_amd_loader_code_object_reader_create_from_file_with_offset_size(
+    hsa_file_t file,
+    size_t offset,
+    size_t size,
+    hsa_code_object_reader_t *code_object_reader) {
+  if (false == core::Runtime::runtime_singleton_->IsOpen()) {
+    return HSA_STATUS_ERROR_NOT_INITIALIZED;
+  }
+  if (nullptr == code_object_reader) {
+    return HSA_STATUS_ERROR_INVALID_ARGUMENT;
+  }
+
+  if (size == 0) {
+    return HSA_STATUS_ERROR_INVALID_CODE_OBJECT;
+  }
+
+  if (__lseek__(file, offset, SEEK_SET) == (off_t)-1) {
+    return HSA_STATUS_ERROR_INVALID_FILE;
+  }
+
+  unsigned char *code_object_memory = new unsigned char[size];
+  if (!code_object_memory) {
+    return HSA_STATUS_ERROR_OUT_OF_RESOURCES;
+  }
+
+  if (__read__(file, code_object_memory, size) != size) {
+    delete [] code_object_memory;
+    return HSA_STATUS_ERROR_INVALID_FILE;
+  }
+
+  CodeObjectReaderWrapper *wrapper = new (std::nothrow) CodeObjectReaderWrapper(
+      code_object_memory, size, offset, file);
+  if (!wrapper) {
+    delete [] code_object_memory;
+    return HSA_STATUS_ERROR_OUT_OF_RESOURCES;
+  }
+
+  *code_object_reader = CodeObjectReaderWrapper::Handle(wrapper);
+  return HSA_STATUS_SUCCESS;
+}

projects/rocr-runtime/runtime/hsa-runtime/inc/hsa_ven_amd_loader.h

@@ -252,10 +252,11 @@ hsa_status_t hsa_ven_amd_loader_query_executable(
  *
  * @param[in] callback Callback to be invoked once per loaded code object. The
  * HSA runtime passes three arguments to the callback: the executable, a
- * loaded code object, and the application data.  If @p callback returns a
+ * loaded code object, and the application data. If @p callback returns a
  * status other than ::HSA_STATUS_SUCCESS for a particular iteration, the
- * traversal stops and ::hsa_executable_iterate_symbols returns that status
- * value.
+ * traversal stops and
+ * ::hsa_ven_amd_loader_executable_iterate_loaded_code_objects returns that
+ * status value.
  *
  * @param[in] data Application data that is passed to @p callback on every
  * iteration. May be NULL.
@@ -343,7 +344,7 @@ typedef enum hsa_ven_amd_loader_loaded_code_object_info_e {
    * attribute is ::int.
    */
   HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_FILE = 7,
-    /**
+  /**
    * The signed byte address difference of the memory address at which the code
    * object is loaded minus the virtual address specified in the code object
    * that is loaded. The value of this attribute is only defined if the
@@ -351,7 +352,7 @@ typedef enum hsa_ven_amd_loader_loaded_code_object_info_e {
    * attribute is ::int64_t.
    */
   HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_DELTA = 8,
-/**
+  /**
    * The base memory address at which the code object is loaded. This is the
    * base address of the allocation for the lowest addressed segment of the code
    * object that is loaded. Note that any non-loaded segments before the first
@@ -365,7 +366,53 @@ typedef enum hsa_ven_amd_loader_loaded_code_object_info_e {
    * value of this attribute is only defined if the executable in which the code
    * object is loaded is froozen. The type of this attribute is ::uint64_t.
    */
-  HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_SIZE = 10
+  HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_SIZE = 10,
+  /**
+   * The length of the URI in bytes, not including the NUL terminator. The type
+   * of this attribute is uint32_t.
+   */
+  HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_URI_LENGTH = 11,
+  /**
+   * The URI name from which the code object was loaded. The type of this
+   * attribute is a NUL terminated \p char* with the length equal to the value
+   * of ::HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_URI_LENGTH attribute.
+   * The URI name syntax is defined by the following BNF syntax:
+   *
+   *     code_object_uri ::== file_uri | memory_uri
+   *     file_uri        ::== "file://" file_path [ range_specifier ]
+   *     memory_uri      ::== "memory://" process_id range_specifier
+   *     range_specifier ::== [ "#" | "?" ] "offset=" number "&" "size=" number
+   *     file_path       ::== URI_ENCODED_OS_FILE_PATH
+   *     process_id      ::== DECIMAL_NUMBER
+   *     number          ::== HEX_NUMBER | DECIMAL_NUMBER | OCTAL_NUMBER
+   *
+   * ``number`` is a C integral literal where hexadecimal values are prefixed by
+   * "0x" or "0X", and octal values by "0".
+   *
+   * ``file_path`` is the file's path specified as a URI encoded UTF-8 string.
+   * In URI encoding, every character that is not in the regular expression
+   * ``[a-zA-Z0-9/_.~-]`` is encoded as two uppercase hexidecimal digits
+   * proceeded by "%".  Directories in the path are separated by "/".
+   *
+   * ``offset`` is a 0-based byte offset to the start of the code object.  For a
+   * file URI, it is from the start of the file specified by the ``file_path``,
+   * and if omitted defaults to 0. For a memory URI, it is the memory address
+   * and is required.
+   *
+   * ``size`` is the number of bytes in the code object.  For a file URI, if
+   * omitted it defaults to the size of the file.  It is required for a memory
+   * URI.
+   *
+   * ``process_id`` is the identity of the process owning the memory.  For Linux
+   * it is the C unsigned integral decimal literal for the process ID (PID).
+   *
+   * For example:
+   *
+   *     file:///dir1/dir2/file1
+   *     file:///dir3/dir4/file2#offset=0x2000&size=3000
+   *     memory://1234#offset=0x20000&size=3000
+   */
+  HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_URI = 12,
 } hsa_ven_amd_loader_loaded_code_object_info_t;
 
 /**
@@ -398,10 +445,58 @@ hsa_status_t hsa_ven_amd_loader_loaded_code_object_get_info(
 
 //===----------------------------------------------------------------------===//
 
+/**
+ * @brief Create a code object reader to operate on a file with size and offset.
+ *
+ * @param[in] file File descriptor. The file must have been opened by
+ * application with at least read permissions prior calling this function. The
+ * file must contain a vendor-specific code object.
+ *
+ * The file is owned and managed by the application; the lifetime of the file
+ * descriptor must exceed that of any associated code object reader.
+ *
+ * @param[in] size Size of the code object embedded in @p file.
+ *
+ * @param[in] offset 0-based offset relative to the beginning of the @p file
+ * that denotes the beginning of the code object embedded within the @p file.
+ *
+ * @param[out] code_object_reader Memory location to store the newly created
+ * code object reader handle. Must not be NULL.
+ *
+ * @retval ::HSA_STATUS_SUCCESS The function has been executed successfully.
+ *
+ * @retval ::HSA_STATUS_ERROR_NOT_INITIALIZED The HSA runtime has not been
+ * initialized.
+ *
+ * @retval ::HSA_STATUS_ERROR_INVALID_FILE @p file is not opened with at least
+ * read permissions. This condition may also be reported as
+ * ::HSA_STATUS_ERROR_INVALID_CODE_OBJECT_READER by the
+ * ::hsa_executable_load_agent_code_object function.
+ *
+ * @retval ::HSA_STATUS_ERROR_INVALID_CODE_OBJECT The bytes starting at offset
+ * do not form a valid code object. If file size is 0. Or offset > file size.
+ * This condition may also be reported as
+ * ::HSA_STATUS_ERROR_INVALID_CODE_OBJECT by the
+ * ::hsa_executable_load_agent_code_object function.
+ *
+ * @retval ::HSA_STATUS_ERROR_OUT_OF_RESOURCES The HSA runtime failed to
+ * allocate the required resources.
+ *
+ * @retval ::HSA_STATUS_ERROR_INVALID_ARGUMENT @p code_object_reader is NULL.
+ */
+hsa_status_t
+hsa_ven_amd_loader_code_object_reader_create_from_file_with_offset_size(
+    hsa_file_t file,
+    size_t offset,
+    size_t size,
+    hsa_code_object_reader_t *code_object_reader);
+
+//===----------------------------------------------------------------------===//
+
 /**
  * @brief Extension version.
  */
-#define hsa_ven_amd_loader 001000
+#define hsa_ven_amd_loader 001002
 
 /**
  * @brief Extension function table version 1.00.
@@ -450,6 +545,43 @@ typedef struct hsa_ven_amd_loader_1_01_pfn_s {
     void *value);
 } hsa_ven_amd_loader_1_01_pfn_t;
 
+/**
+ * @brief Extension function table version 1.02.
+ */
+typedef struct hsa_ven_amd_loader_1_02_pfn_s {
+  hsa_status_t (*hsa_ven_amd_loader_query_host_address)(
+    const void *device_address,
+    const void **host_address);
+
+  hsa_status_t (*hsa_ven_amd_loader_query_segment_descriptors)(
+    hsa_ven_amd_loader_segment_descriptor_t *segment_descriptors,
+    size_t *num_segment_descriptors);
+
+  hsa_status_t (*hsa_ven_amd_loader_query_executable)(
+    const void *device_address,
+    hsa_executable_t *executable);
+
+  hsa_status_t (*hsa_ven_amd_loader_executable_iterate_loaded_code_objects)(
+    hsa_executable_t executable,
+    hsa_status_t (*callback)(
+      hsa_executable_t executable,
+      hsa_loaded_code_object_t loaded_code_object,
+      void *data),
+    void *data);
+
+  hsa_status_t (*hsa_ven_amd_loader_loaded_code_object_get_info)(
+    hsa_loaded_code_object_t loaded_code_object,
+    hsa_ven_amd_loader_loaded_code_object_info_t attribute,
+    void *value);
+
+  hsa_status_t
+    (*hsa_ven_amd_loader_code_object_reader_create_from_file_with_offset_size)(
+      hsa_file_t file,
+      size_t offset,
+      size_t size,
+      hsa_code_object_reader_t *code_object_reader);
+} hsa_ven_amd_loader_1_02_pfn_t;
+
 #ifdef __cplusplus
 }
 #endif /* __cplusplus */

projects/rocr-runtime/runtime/hsa-runtime/loader/executable.cpp

@@ -216,7 +216,7 @@ hsa_status_t AmdHsaCodeLoader::FreezeExecutable(Executable *executable, const ch
     return status;
   }
 
-  // Assumeing runtime atomic implements C++ std::memory_order
+  // Assuming runtime atomic implements C++ std::memory_order
   WriterLockGuard<ReaderWriterLock> writer_lock(rw_lock_);
   atomic::Store(&_amdgpu_r_debug.r_state, r_debug::RT_ADD, std::memory_order_relaxed);
   atomic::Fence(std::memory_order_acq_rel);
@@ -232,7 +232,7 @@ hsa_status_t AmdHsaCodeLoader::FreezeExecutable(Executable *executable, const ch
 }
 
 void AmdHsaCodeLoader::DestroyExecutable(Executable *executable) {
-  // Assumeing runtime atomic implements C++ std::memory_order
+  // Assuming runtime atomic implements C++ std::memory_order
   WriterLockGuard<ReaderWriterLock> writer_lock(rw_lock_);
   atomic::Store(&_amdgpu_r_debug.r_state, r_debug::RT_DELETE, std::memory_order_relaxed);
   atomic::Fence(std::memory_order_acq_rel);
@@ -1018,6 +1018,10 @@ int64_t LoadedCodeObjectImpl::getDelta() const {
   return getLoadBase() - loaded_segments.front()->VAddr();
 }
 
+std::string LoadedCodeObjectImpl::getUri() const {
+  return std::string(r_debug_info.l_name);
+}
+
 hsa_executable_t AmdHsaCodeLoader::FindExecutable(uint64_t device_address)
 {
   hsa_executable_t execHandle = {0};
@@ -1110,9 +1114,10 @@ hsa_status_t ExecutableImpl::LoadCodeObject(
   hsa_agent_t agent,
   hsa_code_object_t code_object,
   const char *options,
+  const std::string &uri,
   hsa_loaded_code_object_t *loaded_code_object)
 {
-  return LoadCodeObject(agent, code_object, 0, options, loaded_code_object);
+  return LoadCodeObject(agent, code_object, 0, options, uri, loaded_code_object);
 }
 
 hsa_status_t ExecutableImpl::LoadCodeObject(
@@ -1120,6 +1125,7 @@ hsa_status_t ExecutableImpl::LoadCodeObject(
   hsa_code_object_t code_object,
   size_t code_object_size,
   const char *options,
+  const std::string &uri,
   hsa_loaded_code_object_t *loaded_code_object)
 {
   WriterLockGuard<ReaderWriterLock> writer_lock(rw_lock_);
@@ -1271,6 +1277,11 @@ hsa_status_t ExecutableImpl::LoadCodeObject(
     }
   }
 
+  loaded_code_objects.back()->r_debug_info.l_addr = loaded_code_objects.back()->getDelta();
+  loaded_code_objects.back()->r_debug_info.l_name = strdup(uri.c_str());
+  loaded_code_objects.back()->r_debug_info.l_prev = nullptr;
+  loaded_code_objects.back()->r_debug_info.l_next = nullptr;
+
   if (nullptr != loaded_code_object) { *loaded_code_object = LoadedCodeObject::Handle(loaded_code_objects.back()); }
   return HSA_STATUS_SUCCESS;
 }
@@ -1862,73 +1873,6 @@ hsa_status_t ExecutableImpl::Freeze(const char *options) {
     for (auto &ls : lco->LoadedSegments()) {
       ls->Freeze();
     }
-    // Update code object debug info after it is frozen.
-    std::stringstream ss;
-    uint64_t elf_begin = lco->getElfData();
-    uint64_t elf_size = lco->getElfSize();
-
-    struct args {
-      ElfW(Addr) mem_addr;
-      size_t callback_num;
-      const char *file_name;
-      size_t file_offset;
-    } data{ elf_begin, 0 };
-
-    // Iterate the loaded shared objects program headers to see if the elf binary
-    // is allocated in a mapped file.
-    if (dl_iterate_phdr([](struct dl_phdr_info *info, size_t size, void *ptr) -> int {
-      struct args *data = (struct args *) ptr;
-      const ElfW(Addr) reladdr = data->mem_addr - info->dlpi_addr;
-
-      int n = info->dlpi_phnum;
-      while (--n >= 0) {
-        if (info->dlpi_phdr[n].p_type == PT_LOAD
-            && reladdr - info->dlpi_phdr[n].p_vaddr >= 0
-            && reladdr - info->dlpi_phdr[n].p_vaddr < info->dlpi_phdr[n].p_memsz) {
-          // The first callback is always the program executable.
-          if (!info->dlpi_name[0] && data->callback_num == 0) {
-            static char argv0[PATH_MAX] = {0};
-            if (!argv0[0] && readlink("/proc/self/exe", argv0, sizeof(argv0)) == -1)
-              return 0;
-            data->file_name = argv0;
-          } else {
-            data->file_name = info->dlpi_name;
-          }
-
-          data->file_offset = reladdr - info->dlpi_phdr[n].p_vaddr + info->dlpi_phdr[n].p_offset;
-          return 1;
-        }
-      }
-
-      ++data->callback_num;
-      return 0;
-    }, &data)) {
-      unsigned char c;
-
-      ss.fill('0');
-      ss << "file://";
-
-      while ((c = *data.file_name++) != '\0') {
-        // %-encode the file name
-        if (isalnum(c) || c == '/' || c == '-' || c == '_' || c == '.' || c == '~') {
-          ss << c;
-        } else {
-          ss << std::uppercase;
-          ss << '%' << std::hex << std::setw(2) << static_cast<int>(c);
-          ss << std::nouppercase;
-        }
-      }
-      ss << "#offset=" << std::dec << data.file_offset
-         << "&size=" << std::dec << elf_size;
-    } else {
-      ss << "file:///proc/" << getpid() << "/mem#"
-         << "offset=" << std::hex << std::showbase << elf_begin << "&"
-         << "size=" << std::dec << elf_size;
-    }
-    lco->r_debug_info.l_addr = lco->getDelta();
-    lco->r_debug_info.l_name = strdup(ss.str().c_str());
-    lco->r_debug_info.l_prev = nullptr;
-    lco->r_debug_info.l_next = nullptr;
   }
 
   state_ = HSA_EXECUTABLE_STATE_FROZEN;

projects/rocr-runtime/runtime/hsa-runtime/loader/executable.hpp

@@ -324,6 +324,7 @@ public:
   uint64_t getLoadBase() const override;
   uint64_t getLoadSize() const override;
   int64_t getDelta() const override;
+  std::string getUri() const override;
 
   link_map r_debug_info;
 };
@@ -435,6 +436,7 @@ public:
     hsa_agent_t agent,
     hsa_code_object_t code_object,
     const char *options,
+    const std::string &uri,
     hsa_loaded_code_object_t *loaded_code_object) override;
 
   hsa_status_t LoadCodeObject(
@@ -442,6 +444,7 @@ public:
     hsa_code_object_t code_object,
     size_t code_object_size,
     const char *options,
+    const std::string &uri,
     hsa_loaded_code_object_t *loaded_code_object) override;
 
   hsa_status_t Freeze(const char *options) override;