diff --git a/runtime/hsa-runtime/core/inc/amd_hsa_loader.hpp b/runtime/hsa-runtime/core/inc/amd_hsa_loader.hpp index 562456a63a..364cad54c3 100644 --- a/runtime/hsa-runtime/core/inc/amd_hsa_loader.hpp +++ b/runtime/hsa-runtime/core/inc/amd_hsa_loader.hpp @@ -74,9 +74,9 @@ typedef hsa_executable_symbol_t hsa_symbol_t; typedef hsa_executable_symbol_info_t hsa_symbol_info_t; /// @brief Loaded code object attributes. -enum hsa_loaded_code_object_info_t { - HSA_LOADED_CODE_OBJECT_INFO_ELF_IMAGE = 0, - HSA_LOADED_CODE_OBJECT_INFO_ELF_IMAGE_SIZE = 1 +enum amd_loaded_code_object_info_t { + AMD_LOADED_CODE_OBJECT_INFO_ELF_IMAGE = 0, + AMD_LOADED_CODE_OBJECT_INFO_ELF_IMAGE_SIZE = 1 }; /// @brief Loaded segment handle. @@ -200,7 +200,7 @@ public: virtual ~LoadedCodeObject() {} - virtual bool GetInfo(hsa_loaded_code_object_info_t attribute, void *value) = 0; + virtual bool GetInfo(amd_loaded_code_object_info_t attribute, void *value) = 0; virtual hsa_status_t IterateLoadedSegments( hsa_status_t (*callback)( @@ -282,16 +282,14 @@ public: hsa_agent_t agent, hsa_code_object_t code_object, const char *options, - hsa_loaded_code_object_t *loaded_code_object = nullptr, - bool load_legacy = true) = 0; + hsa_loaded_code_object_t *loaded_code_object = nullptr) = 0; virtual hsa_status_t LoadCodeObject( hsa_agent_t agent, hsa_code_object_t code_object, size_t code_object_size, const char *options, - hsa_loaded_code_object_t *loaded_code_object = nullptr, - bool load_legacy = true) = 0; + hsa_loaded_code_object_t *loaded_code_object = nullptr) = 0; virtual hsa_status_t Freeze(const char *options) = 0; diff --git a/runtime/hsa-runtime/core/runtime/hsa.cpp b/runtime/hsa-runtime/core/runtime/hsa.cpp index e6711b04d3..cd14411c6c 100644 --- a/runtime/hsa-runtime/core/runtime/hsa.cpp +++ b/runtime/hsa-runtime/core/runtime/hsa.cpp @@ -1926,7 +1926,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, false); + {0}, code_object, options, loaded_code_object); } hsa_status_t hsa_executable_load_agent_code_object( @@ -1951,7 +1951,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, false); + agent, code_object, options, loaded_code_object); } hsa_status_t hsa_executable_freeze( diff --git a/runtime/hsa-runtime/loader/executable.cpp b/runtime/hsa-runtime/loader/executable.cpp index 69a7313e34..b154790020 100644 --- a/runtime/hsa-runtime/loader/executable.cpp +++ b/runtime/hsa-runtime/loader/executable.cpp @@ -529,15 +529,15 @@ bool VariableSymbol::GetInfo(hsa_symbol_info32_t symbol_info, void *value) { return true; } -bool LoadedCodeObjectImpl::GetInfo(hsa_loaded_code_object_info_t attribute, void *value) +bool LoadedCodeObjectImpl::GetInfo(amd_loaded_code_object_info_t attribute, void *value) { assert(value); switch (attribute) { - case HSA_LOADED_CODE_OBJECT_INFO_ELF_IMAGE: + case AMD_LOADED_CODE_OBJECT_INFO_ELF_IMAGE: ((hsa_code_object_t*)value)->handle = reinterpret_cast(elf_data); break; - case HSA_LOADED_CODE_OBJECT_INFO_ELF_IMAGE_SIZE: + case AMD_LOADED_CODE_OBJECT_INFO_ELF_IMAGE_SIZE: *((size_t*)value) = elf_size; break; default: { @@ -988,10 +988,9 @@ hsa_status_t ExecutableImpl::LoadCodeObject( hsa_agent_t agent, hsa_code_object_t code_object, const char *options, - hsa_loaded_code_object_t *loaded_code_object, - bool load_legacy) + hsa_loaded_code_object_t *loaded_code_object) { - return LoadCodeObject(agent, code_object, 0, options, loaded_code_object, load_legacy); + return LoadCodeObject(agent, code_object, 0, options, loaded_code_object); } hsa_status_t ExecutableImpl::LoadCodeObject( @@ -999,8 +998,7 @@ hsa_status_t ExecutableImpl::LoadCodeObject( hsa_code_object_t code_object, size_t code_object_size, const char *options, - hsa_loaded_code_object_t *loaded_code_object, - bool load_legacy) + hsa_loaded_code_object_t *loaded_code_object) { WriterLockGuard writer_lock(rw_lock_); if (HSA_EXECUTABLE_STATE_FROZEN == state_) { @@ -1018,7 +1016,6 @@ hsa_status_t ExecutableImpl::LoadCodeObject( } code.reset(new code::AmdHsaCode()); - load_legacy_ = load_legacy; if (!code->InitAsHandle(code_object)) { return HSA_STATUS_ERROR_INVALID_CODE_OBJECT; @@ -1079,7 +1076,7 @@ hsa_status_t ExecutableImpl::LoadCodeObject( if (status != HSA_STATUS_SUCCESS) return status; for (size_t i = 0; i < code->SymbolCount(); ++i) { - status = LoadSymbol(agent, code->GetSymbol(i)); + status = LoadSymbol(agent, code->GetSymbol(i), majorVersion); if (status != HSA_STATUS_SUCCESS) { return status; } } @@ -1185,20 +1182,24 @@ hsa_status_t ExecutableImpl::LoadSegmentV2(const code::Segment *data_segment, return HSA_STATUS_SUCCESS; } -hsa_status_t ExecutableImpl::LoadSymbol(hsa_agent_t agent, code::Symbol* sym) +hsa_status_t ExecutableImpl::LoadSymbol(hsa_agent_t agent, + code::Symbol* sym, + uint32_t majorVersion) { if (sym->IsDeclaration()) { - return LoadDeclarationSymbol(agent, sym); + return LoadDeclarationSymbol(agent, sym, majorVersion); } else { - return LoadDefinitionSymbol(agent, sym); + return LoadDefinitionSymbol(agent, sym, majorVersion); } } -hsa_status_t ExecutableImpl::LoadDefinitionSymbol(hsa_agent_t agent, code::Symbol* sym) +hsa_status_t ExecutableImpl::LoadDefinitionSymbol(hsa_agent_t agent, + code::Symbol* sym, + uint32_t majorVersion) { bool isAgent = sym->IsAgent(); - if (!load_legacy_) { - isAgent = agent.handle == 0 ? false : true; + if (majorVersion >= 2) { + isAgent = agent.handle != 0; } if (isAgent) { auto agent_symbol = agent_symbols_.find(std::make_pair(sym->Name(), agent)); @@ -1289,7 +1290,9 @@ hsa_status_t ExecutableImpl::LoadDefinitionSymbol(hsa_agent_t agent, code::Symbo return HSA_STATUS_SUCCESS; } -hsa_status_t ExecutableImpl::LoadDeclarationSymbol(hsa_agent_t agent, code::Symbol* sym) +hsa_status_t ExecutableImpl::LoadDeclarationSymbol(hsa_agent_t agent, + code::Symbol* sym, + uint32_t majorVersion) { auto program_symbol = program_symbols_.find(sym->Name()); if (program_symbol == program_symbols_.end()) { diff --git a/runtime/hsa-runtime/loader/executable.hpp b/runtime/hsa-runtime/loader/executable.hpp index 8883fff051..79fcdabfd5 100644 --- a/runtime/hsa-runtime/loader/executable.hpp +++ b/runtime/hsa-runtime/loader/executable.hpp @@ -257,7 +257,7 @@ public: size_t ElfSize() const { return elf_size; } std::vector& LoadedSegments() { return loaded_segments; } - bool GetInfo(hsa_loaded_code_object_info_t attribute, void *value) override; + bool GetInfo(amd_loaded_code_object_info_t attribute, void *value) override; hsa_status_t IterateLoadedSegments( hsa_status_t (*callback)( @@ -376,16 +376,14 @@ public: hsa_agent_t agent, hsa_code_object_t code_object, const char *options, - hsa_loaded_code_object_t *loaded_code_object, - bool load_legacy = true) override; + hsa_loaded_code_object_t *loaded_code_object) override; hsa_status_t LoadCodeObject( hsa_agent_t agent, hsa_code_object_t code_object, size_t code_object_size, const char *options, - hsa_loaded_code_object_t *loaded_code_object, - bool load_legacy = true) override; + hsa_loaded_code_object_t *loaded_code_object) override; hsa_status_t Freeze(const char *options) override; @@ -452,7 +450,6 @@ private: ExecutableImpl& operator=(const ExecutableImpl &e); std::unique_ptr code; - bool load_legacy_; Symbol* GetSymbolInternal( const char *symbol_name, @@ -466,9 +463,9 @@ private: hsa_status_t LoadSegmentV2(const code::Segment *data_segment, loader::Segment *load_segment); - hsa_status_t LoadSymbol(hsa_agent_t agent, amd::hsa::code::Symbol* sym); - hsa_status_t LoadDefinitionSymbol(hsa_agent_t agent, amd::hsa::code::Symbol* sym); - hsa_status_t LoadDeclarationSymbol(hsa_agent_t agent, amd::hsa::code::Symbol* sym); + hsa_status_t LoadSymbol(hsa_agent_t agent, amd::hsa::code::Symbol* sym, uint32_t majorVersion); + hsa_status_t LoadDefinitionSymbol(hsa_agent_t agent, amd::hsa::code::Symbol* sym, uint32_t majorVersion); + hsa_status_t LoadDeclarationSymbol(hsa_agent_t agent, amd::hsa::code::Symbol* sym, uint32_t majorVersion); hsa_status_t ApplyRelocations(hsa_agent_t agent, amd::hsa::code::AmdHsaCode *c); hsa_status_t ApplyStaticRelocationSection(hsa_agent_t agent, amd::hsa::code::RelocationSection* sec);