AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Code object reader improvements

Browse Source 
- 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
This commit is contained in:
Tony
2020-06-20 06:45:32 -04:00
committed by Tony Tye
parent 71d85855d7
commit 91cb98dab6
5 changed files with 486 additions and 326 deletions
Download Patch File Download Diff File Expand all files Collapse all files
runtime/hsa-runtime/core/inc/amd_hsa_loader.hpp
+24 -36
View File
@@ -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<uint64_t>(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<CodeObjectReaderWrapper*>(handle.handle);
CodeObjectReaderImpl *object =
reinterpret_cast<CodeObjectReaderImpl*>(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};
};
//===----------------------------------------------------------------------===//
runtime/hsa-runtime/core/runtime/amd_hsa_loader.cpp
+255 -80
View File
@@ -41,8 +41,13 @@
////////////////////////////////////////////////////////////////////////////////
#include "core/inc/amd_hsa_loader.hpp"
#include "core/inc/runtime.h"
#include <assert.h>
#include <link.h>
#include <linux/limits.h>
#include <sys/mman.h>
#include <stdlib.h>
#include <unistd.h>
#include <cstring>
@@ -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<uintptr_t>(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<uintptr_t>(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<uintptr_t>(code_object_memory);
uintptr_t adjusted_address = address & PAGE_SIZE_MASK;
size_t adjusted_size = code_object_size + (address - adjusted_address);
munmap(reinterpret_cast<void *>(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<unsigned char*>(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<uint64_t>(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<unsigned char> 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
runtime/hsa-runtime/core/runtime/hsa.cpp
+34 -50
View File
@@ -93,7 +93,7 @@ template <class T> struct ValidityError<const T*> {
#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 <class T> struct ValidityError<const T*> {
if (((ptr) == NULL) || !((ptr)->IsValid())) \
return hsa_status_t(ValidityError<decltype(ptr)>::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<CodeObjectReaderImpl> 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<CodeObjectReaderImpl> 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<uint64_t>(wrapper->code_object_memory)};
{reinterpret_cast<uint64_t>(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<uint64_t>(wrapper->code_object_memory)};
return exec->LoadCodeObject(
agent, code_object, options, wrapper->GetUri(), loaded_code_object);
{reinterpret_cast<uint64_t>(reader->GetCodeObjectMemory())};
return exec->LoadCodeObject( agent, code_object, options,
reader->GetUri(), loaded_code_object);
CATCH;
}
runtime/hsa-runtime/core/runtime/hsa_ven_amd_loader.cpp
+167 -160
View File
@@ -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<uint64_t>(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<uintptr_t>(device_address);
uintptr_t uhaddr = Runtime::runtime_singleton_->loader()->FindHostAddress(udaddr);
if (0 == uhaddr) {
return HSA_STATUS_ERROR_INVALID_ARGUMENT;
}
*host_address = reinterpret_cast<void*>(uhaddr);
return HSA_STATUS_SUCCESS;
*host_address = reinterpret_cast<void*>(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<uintptr_t>(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<uint64_t>(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<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;
}
*((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<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;
}
}
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<CodeObjectReaderImpl> 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
runtime/hsa-runtime/core/util/flag.h
+6
View File
@@ -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_;