From 91cb98dab62e498839ec96a990bcf48c0a2862bd Mon Sep 17 00:00:00 2001 From: Tony Date: Sat, 20 Jun 2020 06:45:32 -0400 Subject: [PATCH] Code object reader improvements - Make code object reader use mmap when loading from a file on Linux. - Support computing code object URI for memory either fro the loaded host executables, or from all mmapped files. Define the environment variable HSA_LOADER_ENABLE_MMAP_URI to non 0 to search the mmap files, otherwise only the loaded executables will be seatched. - For mmap search, determine file size and ommit offset and size URI fragment when the code object is the whole file even when specifying a file size explicitly or specifying memory that has been mmaped. - Always return a non-empty code object URI. - When a code object reader is created, complete all fields to ensure it can be used in a multi-threaded manner using only const operations. - Add missing exception handlers in the AMD vendor extentions. - More rigorous checking for errors. Change-Id: I07797b1dc60c5c64245142d77becf9f7c9643395 --- .../hsa-runtime/core/inc/amd_hsa_loader.hpp | 60 ++-- .../core/runtime/amd_hsa_loader.cpp | 335 +++++++++++++----- runtime/hsa-runtime/core/runtime/hsa.cpp | 84 ++--- .../core/runtime/hsa_ven_amd_loader.cpp | 327 ++++++++--------- runtime/hsa-runtime/core/util/flag.h | 6 + 5 files changed, 486 insertions(+), 326 deletions(-) diff --git a/runtime/hsa-runtime/core/inc/amd_hsa_loader.hpp b/runtime/hsa-runtime/core/inc/amd_hsa_loader.hpp index ef356f0a51..f30047d0bf 100644 --- a/runtime/hsa-runtime/core/inc/amd_hsa_loader.hpp +++ b/runtime/hsa-runtime/core/inc/amd_hsa_loader.hpp @@ -107,61 +107,49 @@ namespace amd { namespace hsa { namespace loader { -/// @class CodeObjectReaderWrapper. +/// @class CodeObjectReaderImpl. /// @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; - +struct CodeObjectReaderImpl final { public: /// @returns Handle equivalent of @p object. static hsa_code_object_reader_t Handle( - const CodeObjectReaderWrapper *object) { + const CodeObjectReaderImpl *object) { hsa_code_object_reader_t handle = {reinterpret_cast(object)}; return handle; } /// @returns Object equivalent of @p handle. - static CodeObjectReaderWrapper *Object( + static CodeObjectReaderImpl *Object( const hsa_code_object_reader_t &handle) { - CodeObjectReaderWrapper *object = - reinterpret_cast(handle.handle); + CodeObjectReaderImpl *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, - bool _is_complete_file = false) - : 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) - , is_complete_file(_is_complete_file) {} + CodeObjectReaderImpl() {} /// @brief Default destructor. - ~CodeObjectReaderWrapper() {} + ~CodeObjectReaderImpl(); - bool ComesFromFile() { - return code_object_file_descriptor != -1; - } + hsa_status_t SetFile( + hsa_file_t _code_object_file_descriptor, + size_t _code_object_offset = 0, + size_t _code_object_size = 0); - 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); - } - } + hsa_status_t SetMemory( + const void *_code_object_memory, + size_t _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; - bool is_complete_file; + const void *GetCodeObjectMemory() const { return code_object_memory; }; + + std::string GetUri() const { return uri; }; + + private: + const void *code_object_memory{nullptr}; + size_t code_object_size{0}; + std::string uri{}; + bool is_mmap{false}; }; //===----------------------------------------------------------------------===// diff --git a/runtime/hsa-runtime/core/runtime/amd_hsa_loader.cpp b/runtime/hsa-runtime/core/runtime/amd_hsa_loader.cpp index d081726cf6..c0440c8ee5 100644 --- a/runtime/hsa-runtime/core/runtime/amd_hsa_loader.cpp +++ b/runtime/hsa-runtime/core/runtime/amd_hsa_loader.cpp @@ -41,8 +41,13 @@ //////////////////////////////////////////////////////////////////////////////// #include "core/inc/amd_hsa_loader.hpp" +#include "core/inc/runtime.h" +#include +#include #include +#include +#include #include #include @@ -53,14 +58,26 @@ namespace { -std::string EncodePathname(const char *Pathname) { +#if !defined(_WIN32) && !defined(_WIN64) +uintptr_t PAGE_SIZE_MASK{ + [] () { + uintptr_t page_size = sysconf(_SC_PAGE_SIZE); + if (page_size == -1) { + page_size = 1 << 12; // Default page size to 4KiB. + } + return ~(page_size - 1); + } () + }; +#endif + +std::string EncodePathname(const char *file_path) { std::ostringstream ss; unsigned char c; ss.fill('0'); ss << "file://"; - while ((c = *Pathname++) != '\0') { + while ((c = *file_path++) != '\0') { if (isalnum(c) || c == '/' || c == '-' || c == '_' || c == '.' || c == '~') { ss << c; @@ -74,6 +91,160 @@ std::string EncodePathname(const char *Pathname) { return ss.str(); } +std::string GetUriFromMemoryAddress(const void *memory, size_t size) { + pid_t pid = getpid(); + std::ostringstream uri_stream; + uri_stream << "memory://" << pid + << "#offset=0x" << std::hex << (uintptr_t)memory << std::dec + << "&size=" << size; + return uri_stream.str(); +} + +std::string GetUriFromMemoryInExecutableFile(const void *memory, size_t size) { +#if !defined(_WIN32) && !defined(_WIN64) + uintptr_t address = reinterpret_cast(memory); + struct callback_data_s { + ElfW(Addr) address; + size_t callback_num; + const char *file_path; + size_t file_offset; + } callback_data{address, 0, nullptr, 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 callback_data_s *callback_data = (struct callback_data_s *) ptr; + const ElfW(Addr) elf_address = callback_data->address - info->dlpi_addr; + + int n = info->dlpi_phnum; + while (--n >= 0) { + if (info->dlpi_phdr[n].p_type == PT_LOAD + && elf_address - info->dlpi_phdr[n].p_vaddr >= 0 + && elf_address - 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] && callback_data->callback_num == 0) { + static char argv0[PATH_MAX] = {0}; + if (!argv0[0] && readlink("/proc/self/exe", argv0, sizeof(argv0)) == -1) + return 0; + callback_data->file_path = argv0; + } else { + callback_data->file_path = info->dlpi_name; + } + + callback_data->file_offset = + elf_address - info->dlpi_phdr[n].p_vaddr + info->dlpi_phdr[n].p_offset; + return 1; + } + } + + ++callback_data->callback_num; + return 0; + }, &callback_data)) { + if (!callback_data.file_path || callback_data.file_path[0] == '\0') { + return GetUriFromMemoryAddress(memory, size); + } + + std::ostringstream uri_stream; + uri_stream << EncodePathname(callback_data.file_path); + uri_stream << "#offset=" << callback_data.file_offset; + uri_stream << "&size=" << size; + return uri_stream.str(); + } +#endif // !defined(_WIN32) && !defined(_WIN64) + return GetUriFromMemoryAddress(memory, size); +} + +std::string GetUriFromMemoryInMmapedFile(const void *memory, size_t size) { +#if !defined(_WIN32) && !defined(_WIN64) + std::ifstream proc_maps; + proc_maps.open("/proc/self/maps", std::ifstream::in); + if (!proc_maps.is_open() || !proc_maps.good()) { + return GetUriFromMemoryAddress(memory, size); + } + + std::string line; + while (std::getline(proc_maps, line)) { + std::stringstream tokens(line); + + uintptr_t low_address, high_address; + char dash; + tokens >> std::hex >> low_address >> std::dec + >> dash + >> std::hex >> high_address >> std::dec; + if (dash != '-') { + continue; + } + + uintptr_t address = reinterpret_cast(memory); + if (!(address >= low_address && (address + size) <= high_address)) { + continue; + } + + std::string permissions, device, uri_file_path; + size_t offset; + uint64_t inode; + tokens >> permissions + >> std::hex >> offset >> std::dec + >> device + >> inode + >> uri_file_path; + + if (inode == 0 || uri_file_path.empty()) { + return GetUriFromMemoryAddress(memory, size); + } + + size_t uri_offset = offset + address - low_address; + + bool is_complete_file = false; + if (uri_offset == 0) { + std::ifstream uri_file(uri_file_path, std::ios::binary); + if (uri_file) { + uri_file.seekg(0, std::ios::end); + is_complete_file = uri_file.tellg() == size; + } + } + + std::ostringstream uri_stream; + uri_stream << EncodePathname(uri_file_path.c_str()); + if (!is_complete_file) { + uri_stream << "#offset=" << uri_offset; + uri_stream << "&size=" << size; + } + return uri_stream.str(); + } +#endif // !defined(_WIN32) && !defined(_WIN64) + return GetUriFromMemoryAddress(memory, size); +} + +std::string GetUriFromFile(int file_descriptor, size_t offset, size_t size, + bool is_complete_file, const void *memory) { +#if !defined(_WIN32) && !defined(_WIN64) + std::ostringstream proc_fd_path; + proc_fd_path << "/proc/self/fd/" << file_descriptor; + + char uri_file_path[PATH_MAX]; + memset(uri_file_path, 0, PATH_MAX); + + if (readlink(proc_fd_path.str().c_str(), uri_file_path, PATH_MAX) == -1) { + return GetUriFromMemoryAddress(memory, size); + } + + if (uri_file_path[0] == '\0') { + return GetUriFromMemoryAddress(memory, size); + } + + std::ostringstream uri_stream; + uri_stream << EncodePathname(uri_file_path); + if (!is_complete_file) { + uri_stream << "#offset=" << offset; + uri_stream << "&size=" << size; + } + return uri_stream.str(); +#else + return GetUriFromMemoryAddress(memory, size); +#endif // !defined(_WIN32) && !defined(_WIN64) +} + } // namespace namespace rocr { @@ -81,96 +252,100 @@ namespace amd { namespace hsa { namespace loader { -std::string CodeObjectReaderWrapper::GetUriFromFile( - int Fd, size_t Offset, size_t Size) const { +/// @brief Default destructor. +CodeObjectReaderImpl::~CodeObjectReaderImpl() { + if (is_mmap) { #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 (!is_complete_file) { - UriStream << "#offset=" << Offset; - UriStream << "&size=" << Size; - } - return UriStream.str(); + uintptr_t address = reinterpret_cast(code_object_memory); + uintptr_t adjusted_address = address & PAGE_SIZE_MASK; + size_t adjusted_size = code_object_size + (address - adjusted_address); + munmap(reinterpret_cast(adjusted_address), adjusted_size); #else - return std::string(); + delete [] code_object_memory; #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(); -} +hsa_status_t CodeObjectReaderImpl::SetFile( + hsa_file_t _code_object_file_descriptor, + size_t _code_object_offset, + size_t _code_object_size) { + assert(!code_object_memory && "Code object reader wrapper is already set"); + + if (_code_object_file_descriptor == -1) { + return HSA_STATUS_ERROR_INVALID_ARGUMENT; + } + + off_t file_size = __lseek__(_code_object_file_descriptor, 0, SEEK_END); + if (file_size == (off_t)-1) { + return HSA_STATUS_ERROR_INVALID_FILE; + } + if (file_size <= _code_object_offset) { + return HSA_STATUS_ERROR_INVALID_CODE_OBJECT; + } + if (_code_object_size == 0) { + _code_object_size = file_size - _code_object_offset; + } + bool is_complete_file = _code_object_offset == 0 && _code_object_size == file_size; -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); + off_t adjusted_offset = _code_object_offset & PAGE_SIZE_MASK; + size_t adjusted_size = _code_object_size + (_code_object_offset - adjusted_offset); + void *memory = mmap(nullptr, adjusted_size, PROT_READ, MAP_PRIVATE, + _code_object_file_descriptor, adjusted_offset); + if (memory == (void *) -1) { + return HSA_STATUS_ERROR_INVALID_FILE; + } + code_object_memory = reinterpret_cast(memory) + + (_code_object_offset & ~PAGE_SIZE_MASK); + code_object_size = _code_object_size; + is_mmap = true; +#else + if (__lseek__(_code_object_file_descriptor, 0, SEEK_SET) == (off_t)-1) { + return HSA_STATUS_ERROR_INVALID_FILE; } - 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 + Size) <= 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); - } - - uint64_t UriOffset = Offset + MyAddress - LowAddress; - - std::ostringstream UriStream; - UriStream << EncodePathname(Pathname.c_str()); - UriStream << "#offset=" << UriOffset; - if (Size) { - UriStream << "&size=" << Size; - } - return UriStream.str(); + std::unique_ptr memory(new unsigned char[_code_object_size]); + if (!memory) { + return HSA_STATUS_ERROR_OUT_OF_RESOURCES; } + if (__read__(_code_object_file_descriptor, mmap_memory, + _code_object_size) != _code_object_size) { + return HSA_STATUS_ERROR_INVALID_FILE; + } + mmap_memory = memory.release(); + mmap_size = _code_object_size; + code_object_memory = memory; + code_object_size = _code_object_size; #endif // !defined(_WIN32) && !defined(_WIN64) - return GetUriFromMemoryBasic(Mem, Size); + + uri = GetUriFromFile(_code_object_file_descriptor, _code_object_offset, + _code_object_size, is_complete_file, code_object_memory); + + return HSA_STATUS_SUCCESS; +} + +hsa_status_t CodeObjectReaderImpl::SetMemory( + const void *_code_object_memory, + size_t _code_object_size) { + assert(!code_object_memory && "Code object reader wrapper is already set"); + + if (!_code_object_memory || _code_object_size == 0) { + return HSA_STATUS_ERROR_INVALID_ARGUMENT; + } + + code_object_memory = _code_object_memory; + code_object_size = _code_object_size; + + bool loader_enable_mmap_uri = core::Runtime::runtime_singleton_->flag().loader_enable_mmap_uri(); + if (loader_enable_mmap_uri) { + uri = GetUriFromMemoryInMmapedFile(_code_object_memory, _code_object_size); + } else { + uri = GetUriFromMemoryInExecutableFile(_code_object_memory, _code_object_size); + } + + return HSA_STATUS_SUCCESS; } } // namespace loader diff --git a/runtime/hsa-runtime/core/runtime/hsa.cpp b/runtime/hsa-runtime/core/runtime/hsa.cpp index 6cdbc1996c..105960f2da 100644 --- a/runtime/hsa-runtime/core/runtime/hsa.cpp +++ b/runtime/hsa-runtime/core/runtime/hsa.cpp @@ -93,7 +93,7 @@ template struct ValidityError { #define IS_BAD_PTR(ptr) \ do { \ - if ((ptr) == NULL) return HSA_STATUS_ERROR_INVALID_ARGUMENT; \ + if ((ptr) == nullptr) return HSA_STATUS_ERROR_INVALID_ARGUMENT; \ } while (false) #define IS_BAD_PROFILE(profile) \ do { \ @@ -137,9 +137,13 @@ template struct ValidityError { if (((ptr) == NULL) || !((ptr)->IsValid())) \ return hsa_status_t(ValidityError::kValue); \ } while (false) +#define CHECK_STATUS(status) \ + do { \ + if ((status) != HSA_STATUS_SUCCESS) return status; \ + } while (false) #define CHECK_ALLOC(ptr) \ do { \ - if ((ptr) == NULL) return HSA_STATUS_ERROR_OUT_OF_RESOURCES; \ + if ((ptr) == nullptr) return HSA_STATUS_ERROR_OUT_OF_RESOURCES; \ } while (false) #define IS_OPEN() \ do { \ @@ -2045,7 +2049,7 @@ hsa_status_t hsa_code_object_iterate_symbols( using amd::hsa::common::Signed; using amd::hsa::loader::Loader; using amd::hsa::loader::Executable; -using amd::hsa::loader::CodeObjectReaderWrapper; +using amd::hsa::loader::CodeObjectReaderImpl; namespace { @@ -2062,35 +2066,14 @@ hsa_status_t hsa_code_object_reader_create_from_file( IS_OPEN(); IS_BAD_PTR(code_object_reader); - off_t file_size = __lseek__(file, 0, SEEK_END); - if (file_size == (off_t)-1) { - return HSA_STATUS_ERROR_INVALID_FILE; - } + std::unique_ptr reader( + new (std::nothrow) CodeObjectReaderImpl()); + CHECK_ALLOC(reader); - if (file_size == 0) { - return HSA_STATUS_ERROR_INVALID_CODE_OBJECT; - } + hsa_status_t status = reader->SetFile(file); + CHECK_STATUS(status); - if (__lseek__(file, 0, SEEK_SET) == (off_t)-1) { - return HSA_STATUS_ERROR_INVALID_FILE; - } - - unsigned char *code_object_memory = new unsigned char[file_size]; - CHECK_ALLOC(code_object_memory); - - if (__read__(file, code_object_memory, file_size) != file_size) { - delete [] code_object_memory; - return HSA_STATUS_ERROR_INVALID_FILE; - } - - CodeObjectReaderWrapper *wrapper = new (std::nothrow) CodeObjectReaderWrapper( - code_object_memory, file_size, 0, file, true); - if (!wrapper) { - delete [] code_object_memory; - return HSA_STATUS_ERROR_OUT_OF_RESOURCES; - } - - *code_object_reader = CodeObjectReaderWrapper::Handle(wrapper); + *code_object_reader = CodeObjectReaderImpl::Handle(reader.release()); return HSA_STATUS_SUCCESS; CATCH; } @@ -2108,11 +2091,15 @@ hsa_status_t hsa_code_object_reader_create_from_memory( return HSA_STATUS_ERROR_INVALID_ARGUMENT; } - CodeObjectReaderWrapper *wrapper = new (std::nothrow) CodeObjectReaderWrapper( - code_object, size, 0, -1); - CHECK_ALLOC(wrapper); + std::unique_ptr reader( + new (std::nothrow) CodeObjectReaderImpl()); + CHECK_ALLOC(reader); - *code_object_reader = CodeObjectReaderWrapper::Handle(wrapper); + hsa_status_t status = reader->SetMemory(code_object, size); + CHECK_STATUS(status); + + *code_object_reader = + CodeObjectReaderImpl::Handle(reader.release()); return HSA_STATUS_SUCCESS; CATCH; } @@ -2122,16 +2109,13 @@ hsa_status_t hsa_code_object_reader_destroy( TRY; IS_OPEN(); - CodeObjectReaderWrapper *wrapper = CodeObjectReaderWrapper::Object( - code_object_reader); - if (!wrapper) { + CodeObjectReaderImpl *reader = + CodeObjectReaderImpl::Object(code_object_reader); + if (!reader) { return HSA_STATUS_ERROR_INVALID_CODE_OBJECT_READER; } - if (wrapper->ComesFromFile()) { - delete [] (unsigned char*)wrapper->code_object_memory; - } - delete wrapper; + delete reader; return HSA_STATUS_SUCCESS; CATCH; @@ -2238,16 +2222,16 @@ hsa_status_t hsa_executable_load_program_code_object( return HSA_STATUS_ERROR_INVALID_EXECUTABLE; } - CodeObjectReaderWrapper *wrapper = CodeObjectReaderWrapper::Object( + CodeObjectReaderImpl *reader = CodeObjectReaderImpl::Object( code_object_reader); - if (!wrapper) { + if (!reader) { return HSA_STATUS_ERROR_INVALID_CODE_OBJECT_READER; } hsa_code_object_t code_object = - {reinterpret_cast(wrapper->code_object_memory)}; + {reinterpret_cast(reader->GetCodeObjectMemory())}; return exec->LoadCodeObject( - {0}, code_object, options, wrapper->GetUri(), loaded_code_object); + {0}, code_object, options, reader->GetUri(), loaded_code_object); CATCH; } @@ -2265,16 +2249,16 @@ hsa_status_t hsa_executable_load_agent_code_object( return HSA_STATUS_ERROR_INVALID_EXECUTABLE; } - CodeObjectReaderWrapper *wrapper = CodeObjectReaderWrapper::Object( + CodeObjectReaderImpl *reader = CodeObjectReaderImpl::Object( code_object_reader); - if (!wrapper) { + if (!reader) { return HSA_STATUS_ERROR_INVALID_CODE_OBJECT_READER; } hsa_code_object_t code_object = - {reinterpret_cast(wrapper->code_object_memory)}; - return exec->LoadCodeObject( - agent, code_object, options, wrapper->GetUri(), loaded_code_object); + {reinterpret_cast(reader->GetCodeObjectMemory())}; + return exec->LoadCodeObject( agent, code_object, options, + reader->GetUri(), loaded_code_object); CATCH; } diff --git a/runtime/hsa-runtime/core/runtime/hsa_ven_amd_loader.cpp b/runtime/hsa-runtime/core/runtime/hsa_ven_amd_loader.cpp index a806e3f349..a74181d3b3 100644 --- a/runtime/hsa-runtime/core/runtime/hsa_ven_amd_loader.cpp +++ b/runtime/hsa-runtime/core/runtime/hsa_ven_amd_loader.cpp @@ -46,67 +46,78 @@ #include "core/inc/runtime.h" namespace rocr { - + using namespace amd::hsa; using namespace core; -using loader::CodeObjectReaderWrapper; +using loader::CodeObjectReaderImpl; using loader::Executable; using loader::LoadedCodeObject; +namespace AMD { + +hsa_status_t handleException(); + +} // namespace amd + hsa_status_t hsa_ven_amd_loader_query_host_address( const void *device_address, const void **host_address) { - if (false == Runtime::runtime_singleton_->IsOpen()) { - return HSA_STATUS_ERROR_NOT_INITIALIZED; - } - if (nullptr == device_address) { - return HSA_STATUS_ERROR_INVALID_ARGUMENT; - } - if (nullptr == host_address) { - return HSA_STATUS_ERROR_INVALID_ARGUMENT; - } + try { + if (!Runtime::runtime_singleton_->IsOpen()) { + return HSA_STATUS_ERROR_NOT_INITIALIZED; + } + if (nullptr == device_address) { + return HSA_STATUS_ERROR_INVALID_ARGUMENT; + } + if (nullptr == host_address) { + return HSA_STATUS_ERROR_INVALID_ARGUMENT; + } - uint64_t udaddr = reinterpret_cast(device_address); - uint64_t uhaddr = Runtime::runtime_singleton_->loader()->FindHostAddress(udaddr); - if (0 == uhaddr) { - return HSA_STATUS_ERROR_INVALID_ARGUMENT; - } + uintptr_t udaddr = reinterpret_cast(device_address); + uintptr_t uhaddr = Runtime::runtime_singleton_->loader()->FindHostAddress(udaddr); + if (0 == uhaddr) { + return HSA_STATUS_ERROR_INVALID_ARGUMENT; + } - *host_address = reinterpret_cast(uhaddr); - return HSA_STATUS_SUCCESS; + *host_address = reinterpret_cast(uhaddr); + return HSA_STATUS_SUCCESS; + } catch(...) { return AMD::handleException(); } } hsa_status_t hsa_ven_amd_loader_query_segment_descriptors( hsa_ven_amd_loader_segment_descriptor_t *segment_descriptors, size_t *num_segment_descriptors) { - if (false == Runtime::runtime_singleton_->IsOpen()) { - return HSA_STATUS_ERROR_NOT_INITIALIZED; - } + try { + if (!Runtime::runtime_singleton_->IsOpen()) { + return HSA_STATUS_ERROR_NOT_INITIALIZED; + } - // Arguments are checked by the loader. - return Runtime::runtime_singleton_->loader()->QuerySegmentDescriptors(segment_descriptors, num_segment_descriptors); + // Arguments are checked by the loader. + return Runtime::runtime_singleton_->loader()->QuerySegmentDescriptors(segment_descriptors, num_segment_descriptors); + } catch(...) { return AMD::handleException(); } } hsa_status_t hsa_ven_amd_loader_query_executable( const void *device_address, hsa_executable_t *executable) { + try { + if (!Runtime::runtime_singleton_->IsOpen()) { + return HSA_STATUS_ERROR_NOT_INITIALIZED; + } + if ((nullptr == device_address) || (nullptr == executable)) { + return HSA_STATUS_ERROR_INVALID_ARGUMENT; + } - if (false == Runtime::runtime_singleton_->IsOpen()) { - return HSA_STATUS_ERROR_NOT_INITIALIZED; - } - if ((nullptr == device_address) || (nullptr == executable)) { - return HSA_STATUS_ERROR_INVALID_ARGUMENT; - } + uintptr_t udaddr = reinterpret_cast(device_address); + hsa_executable_t exec = Runtime::runtime_singleton_->loader()->FindExecutable(udaddr); + if (0 == exec.handle) { + return HSA_STATUS_ERROR_INVALID_ARGUMENT; + } - uint64_t udaddr = reinterpret_cast(device_address); - hsa_executable_t exec = Runtime::runtime_singleton_->loader()->FindExecutable(udaddr); - if (0 == exec.handle) { - return HSA_STATUS_ERROR_INVALID_ARGUMENT; - } - - *executable = exec; - return HSA_STATUS_SUCCESS; + *executable = exec; + return HSA_STATUS_SUCCESS; + } catch(...) { return AMD::handleException(); } } hsa_status_t hsa_ven_amd_loader_executable_iterate_loaded_code_objects( @@ -116,112 +127,116 @@ hsa_status_t hsa_ven_amd_loader_executable_iterate_loaded_code_objects( hsa_loaded_code_object_t loaded_code_object, void *data), void *data) { - if (false == Runtime::runtime_singleton_->IsOpen()) { - return HSA_STATUS_ERROR_NOT_INITIALIZED; - } - if (nullptr == callback) { - return HSA_STATUS_ERROR_INVALID_ARGUMENT; - } + try { + if (!Runtime::runtime_singleton_->IsOpen()) { + return HSA_STATUS_ERROR_NOT_INITIALIZED; + } + if (nullptr == callback) { + return HSA_STATUS_ERROR_INVALID_ARGUMENT; + } - Executable *exec = Executable::Object(executable); - if (!exec) { - return HSA_STATUS_ERROR_INVALID_EXECUTABLE; - } + Executable *exec = Executable::Object(executable); + if (!exec) { + return HSA_STATUS_ERROR_INVALID_EXECUTABLE; + } - return exec->IterateLoadedCodeObjects(callback, data); + return exec->IterateLoadedCodeObjects(callback, data); + } catch(...) { return AMD::handleException(); } } 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) { - if (false == Runtime::runtime_singleton_->IsOpen()) { - return HSA_STATUS_ERROR_NOT_INITIALIZED; - } - if (nullptr == value) { - return HSA_STATUS_ERROR_INVALID_ARGUMENT; - } - - const LoadedCodeObject *lcobj = LoadedCodeObject::Object(loaded_code_object); - if (!lcobj) { - return HSA_STATUS_ERROR_INVALID_CODE_OBJECT; - } - - switch (attribute) { - case HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_EXECUTABLE: { - *((hsa_executable_t*)value) = lcobj->getExecutable(); - break; + try { + if (!Runtime::runtime_singleton_->IsOpen()) { + return HSA_STATUS_ERROR_NOT_INITIALIZED; } - case HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_KIND: { - *((uint32_t*)value) = lcobj->getAgent().handle == 0 - ? HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_KIND_PROGRAM - : HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_KIND_AGENT; - break; + if (nullptr == value) { + return HSA_STATUS_ERROR_INVALID_ARGUMENT; } - case HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_AGENT: { - hsa_agent_t agent = lcobj->getAgent(); - if (agent.handle == 0) { - return HSA_STATUS_ERROR_INVALID_ARGUMENT; + + const LoadedCodeObject *lcobj = LoadedCodeObject::Object(loaded_code_object); + if (!lcobj) { + return HSA_STATUS_ERROR_INVALID_CODE_OBJECT; + } + + switch (attribute) { + case HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_EXECUTABLE: { + *((hsa_executable_t*)value) = lcobj->getExecutable(); + break; + } + case HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_KIND: { + *((uint32_t*)value) = lcobj->getAgent().handle == 0 + ? HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_KIND_PROGRAM + : HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_KIND_AGENT; + break; + } + case HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_AGENT: { + hsa_agent_t agent = lcobj->getAgent(); + if (agent.handle == 0) { + return HSA_STATUS_ERROR_INVALID_ARGUMENT; + } + *((hsa_agent_t*)value) = agent; + break; + } + case HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_TYPE: { + // TODO Update loader so it keeps track if code object was loaded from a + // file or memory. + *((uint32_t*)value) = HSA_VEN_AMD_LOADER_CODE_OBJECT_STORAGE_TYPE_MEMORY; + break; + } + case HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_MEMORY_BASE: { + *((uint64_t*)value) = lcobj->getElfData(); + break; + } + case HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_MEMORY_SIZE: { + *((uint64_t*)value) = lcobj->getElfSize(); + break; + } + case HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_FILE: { + // TODO Update loader so it keeps track if code object was loaded from a + // file or memory. + return HSA_STATUS_ERROR_INVALID_ARGUMENT; + break; + } + case HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_DELTA: { + // TODO Check if executable is frozen. + // This suggests this code should be moved into LoadedCodeObjectImpl::getinfo + // as is done for other *_get_info methods. Currently LoadedCodeObject has a + // GetInfo method which is likely not used. + // Also should this have a *NOT_FROZEN ststus code added? + // if (state_ != HSA_EXECUTABLE_STATE_FROZEN) { + // return HSA_STATUS_ERROR_INVALID_ARGUMENT; + // } + *((int64_t*)value) = lcobj->getDelta(); + break; + } + case HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_BASE: { + // TODO Check if executable is frozen. + *((uint64_t*)value) = lcobj->getLoadBase(); + break; + } + case HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_SIZE: { + // TODO Check if executable is frozen. + *((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; } - *((hsa_agent_t*)value) = agent; - break; } - case HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_TYPE: { - // TODO Update loader so it keeps track if code object was loaded from a - // file or memory. - *((uint32_t*)value) = HSA_VEN_AMD_LOADER_CODE_OBJECT_STORAGE_TYPE_MEMORY; - break; - } - case HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_MEMORY_BASE: { - *((uint64_t*)value) = lcobj->getElfData(); - break; - } - case HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_MEMORY_SIZE: { - *((uint64_t*)value) = lcobj->getElfSize(); - break; - } - case HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_CODE_OBJECT_STORAGE_FILE: { - // TODO Update loader so it keeps track if code object was loaded from a - // file or memory. - return HSA_STATUS_ERROR_INVALID_ARGUMENT; - break; - } - case HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_DELTA: { - // TODO Check if executable is frozen. - // This suggests this code should be moved into LoadedCodeObjectImpl::getinfo - // as is done for other *_get_info methods. Currently LoadedCodeObject has a - // GetInfo method which is likely not used. - // Also should this have a *NOT_FROZEN ststus code added? - // if (state_ != HSA_EXECUTABLE_STATE_FROZEN) { - // return HSA_STATUS_ERROR_INVALID_ARGUMENT; - // } - *((int64_t*)value) = lcobj->getDelta(); - break; - } - case HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_BASE: { - // TODO Check if executable is frozen. - *((uint64_t*)value) = lcobj->getLoadBase(); - break; - } - case HSA_VEN_AMD_LOADER_LOADED_CODE_OBJECT_INFO_LOAD_SIZE: { - // TODO Check if executable is frozen. - *((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; - } - } - return HSA_STATUS_SUCCESS; + return HSA_STATUS_SUCCESS; + } catch(...) { return AMD::handleException(); } } hsa_status_t @@ -230,40 +245,32 @@ hsa_ven_amd_loader_code_object_reader_create_from_file_with_offset_size( size_t offset, size_t size, hsa_code_object_reader_t *code_object_reader) { - if (false == Runtime::runtime_singleton_->IsOpen()) { - return HSA_STATUS_ERROR_NOT_INITIALIZED; - } - if (nullptr == code_object_reader) { - return HSA_STATUS_ERROR_INVALID_ARGUMENT; - } + try { + if (!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 (size == 0) { + return HSA_STATUS_ERROR_INVALID_CODE_OBJECT; + } - if (__lseek__(file, offset, SEEK_SET) == (off_t)-1) { - return HSA_STATUS_ERROR_INVALID_FILE; - } + std::unique_ptr reader( + new (std::nothrow) CodeObjectReaderImpl()); + if (!reader) { + return HSA_STATUS_ERROR_OUT_OF_RESOURCES; + } - unsigned char *code_object_memory = new unsigned char[size]; - if (!code_object_memory) { - return HSA_STATUS_ERROR_OUT_OF_RESOURCES; - } + hsa_status_t status = reader->SetFile(file, offset, size); + if (status != HSA_STATUS_SUCCESS) { + return status; + } - 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; + *code_object_reader = CodeObjectReaderImpl::Handle(reader.release()); + return HSA_STATUS_SUCCESS; + } catch(...) { return AMD::handleException(); } } } // namespace rocr diff --git a/runtime/hsa-runtime/core/util/flag.h b/runtime/hsa-runtime/core/util/flag.h index e0df01c91b..f43c6da32a 100644 --- a/runtime/hsa-runtime/core/util/flag.h +++ b/runtime/hsa-runtime/core/util/flag.h @@ -118,6 +118,9 @@ class Flag { var = os::GetEnvVar("HSA_DISABLE_IMAGE"); disable_image_ = (var == "1") ? true : false; + + var = os::GetEnvVar("HSA_LOADER_ENABLE_MMAP_URI"); + loader_enable_mmap_uri_ = (var == "1") ? true : false; } bool check_flat_scratch() const { return check_flat_scratch_; } @@ -160,6 +163,8 @@ class Flag { bool disable_image() const { return disable_image_; } + bool loader_enable_mmap_uri() const { return loader_enable_mmap_uri_; } + private: bool check_flat_scratch_; bool enable_vm_fault_message_; @@ -175,6 +180,7 @@ class Flag { bool no_scratch_reclaim_; bool no_scratch_thread_limit_; bool disable_image_; + bool loader_enable_mmap_uri_; std::string enable_sdma_;