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

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;

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

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

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