diff --git a/projects/rocr-runtime/runtime/hsa-runtime/CMakeLists.txt b/projects/rocr-runtime/runtime/hsa-runtime/CMakeLists.txt index 7174044dc3..eab7768217 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/CMakeLists.txt +++ b/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" diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_hsa_loader.hpp b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_hsa_loader.hpp index a7f2d6a961..0073053581 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_hsa_loader.hpp +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_hsa_loader.hpp @@ -53,6 +53,16 @@ #include #include +#if defined(_WIN32) || defined(_WIN64) +#include +#define __read__ _read +#define __lseek__ _lseek +#else +#include +#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(object)}; + return handle; + } + + /// @returns Object equivalent of @p handle. + static CodeObjectReaderWrapper *Object( + const hsa_code_object_reader_t &handle) { + CodeObjectReaderWrapper *object = + reinterpret_cast(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; diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_hsa_loader.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_hsa_loader.cpp new file mode 100644 index 0000000000..9663dfa540 --- /dev/null +++ b/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 +#include + +#include +#include +#include +#include +#include + +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(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(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 diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa.cpp index 1d08c25643..80a3eb7113 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa.cpp +++ b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa.cpp @@ -47,16 +47,6 @@ #include #include -#if defined(_WIN32) || defined(_WIN64) -#include -#define __read__ _read -#define __lseek__ _lseek -#else -#include -#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(object)}; - return handle; - } - - /// @returns Object equivalent of @p handle. - static CodeObjectReaderWrapper *Object( - const hsa_code_object_reader_t &handle) { - CodeObjectReaderWrapper *object = common::ObjectAt( - 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(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(wrapper->code_object_memory)}; return exec->LoadCodeObject( - agent, code_object, options, loaded_code_object); + agent, code_object, options, wrapper->GetUri(), loaded_code_object); CATCH; } diff --git a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa_ven_amd_loader.cpp b/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa_ven_amd_loader.cpp index 7277ea76ad..c911de107e 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/core/runtime/hsa_ven_amd_loader.cpp +++ b/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(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; +} diff --git a/projects/rocr-runtime/runtime/hsa-runtime/inc/hsa_ven_amd_loader.h b/projects/rocr-runtime/runtime/hsa-runtime/inc/hsa_ven_amd_loader.h index 69a6b18a06..93437fc94b 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/inc/hsa_ven_amd_loader.h +++ b/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 */ diff --git a/projects/rocr-runtime/runtime/hsa-runtime/loader/executable.cpp b/projects/rocr-runtime/runtime/hsa-runtime/loader/executable.cpp index 25c30736ac..d2935a3f97 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/loader/executable.cpp +++ b/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 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 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 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(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; diff --git a/projects/rocr-runtime/runtime/hsa-runtime/loader/executable.hpp b/projects/rocr-runtime/runtime/hsa-runtime/loader/executable.hpp index e9318f9924..7e3ae6c200 100644 --- a/projects/rocr-runtime/runtime/hsa-runtime/loader/executable.hpp +++ b/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;