From 5a8c84e012a6fffc4be84ae261ee349f5c89b8ce Mon Sep 17 00:00:00 2001 From: "Nikolay Haustov [TEXT]" Date: Thu, 28 May 2015 05:38:18 -0500 Subject: [PATCH] ECR #010005 - Update HSA samples and test to use libHSAIL high-level tool interface. [git-p4: depot-paths = "//depot/stg/hsa/drivers/hsa/runtime/": change = 1155310] --- samples/common/assemble.cpp | 167 ---------------------------- samples/common/assemble.hpp | 50 --------- samples/common/hsa_base_util.cpp | 7 +- samples/common/hsa_base_util.h | 3 +- samples/common/hsa_rsrc_factory.cpp | 8 +- samples/common/hsa_rsrc_factory.hpp | 2 + 6 files changed, 14 insertions(+), 223 deletions(-) delete mode 100644 samples/common/assemble.cpp delete mode 100644 samples/common/assemble.hpp diff --git a/samples/common/assemble.cpp b/samples/common/assemble.cpp deleted file mode 100644 index 2e1827299c..0000000000 --- a/samples/common/assemble.cpp +++ /dev/null @@ -1,167 +0,0 @@ -#include -#include -#include -#include -#include -#include -#include -#include -#include "assemble.hpp" -#include "hsa.h" -#include "hsa_ext_finalize.h" -#include "HSAILDisassembler.h" -#include "HSAILParser.h" -#include "HSAILScanner.h" -#include "HSAILValidator.h" -#include "HSAILBrigObjectFile.h" - -namespace { - std::unordered_map mod2con; -} // namespace anonymous - -hsa_status_t ModuleCreateFromHsailTextFile( - const char *hsail_text_filename, - hsa_ext_module_t *module -) { - if (!hsail_text_filename) { - return HSA_STATUS_ERROR_INVALID_ARGUMENT; - } - - std::ifstream hsail_text_file(hsail_text_filename); - if (!hsail_text_file.is_open()) { - return HSA_STATUS_ERROR_INVALID_ARGUMENT; - } - - std::string hsail_string; - hsail_text_file.seekg(0, std::ios::end); - hsail_string.resize(hsail_text_file.tellg()); - hsail_text_file.seekg(0, std::ios::beg); - hsail_text_file.read(&hsail_string[0], hsail_string.size()); - hsail_text_file.close(); - - return ModuleCreateFromHsailString(hsail_string.c_str(), module); -} - -hsa_status_t ModuleCreateFromBrigFile( - const char *filename, - hsa_ext_module_t *module -) { - HSAIL_ASM::BrigContainer *brig_container = new HSAIL_ASM::BrigContainer; - std::stringstream ss; - int rc = HSAIL_ASM::BrigIO::load(*brig_container, HSAIL_ASM::FILE_FORMAT_AUTO, HSAIL_ASM::BrigIO::fileReadingAdapter(filename, ss)); - if (rc != 0) { return static_cast(HSA_EXT_STATUS_ERROR_INVALID_MODULE); } - auto insert_status = mod2con.insert( - std::make_pair( - brig_container->getBrigModule(), - reinterpret_cast(brig_container) - ) - ); - assert(insert_status.second); - *module = brig_container->getBrigModule(); - return HSA_STATUS_SUCCESS; -} - -hsa_status_t ModuleCreateFromHsailString( - const char *hsail_string, - hsa_ext_module_t *module -) { - if (!hsail_string || !module) { - return HSA_STATUS_ERROR_INVALID_ARGUMENT; - } - - HSAIL_ASM::BrigContainer *brig_container = NULL; - try { - brig_container = new HSAIL_ASM::BrigContainer; - } catch (const std::bad_alloc) { - return HSA_STATUS_ERROR_OUT_OF_RESOURCES; - } - - std::istringstream hsail_stream(hsail_string); - try { - HSAIL_ASM::Scanner hsail_scanner(hsail_stream); - HSAIL_ASM::Parser hsail_parser(hsail_scanner, *brig_container); - hsail_parser.parseSource(); - } catch (const SyntaxError) { - delete brig_container; - return static_cast(HSA_EXT_STATUS_ERROR_INVALID_MODULE); - } - - try { - auto insert_status = mod2con.insert( - std::make_pair( - brig_container->getBrigModule(), - reinterpret_cast(brig_container) - ) - ); - assert(insert_status.second); - } catch (const std::bad_alloc) { - delete brig_container; - return HSA_STATUS_ERROR_OUT_OF_RESOURCES; - } - - *module = brig_container->getBrigModule(); - return HSA_STATUS_SUCCESS; -} - -hsa_status_t ModuleDestroy( - hsa_ext_module_t module -) { - auto find_status = mod2con.find(module); - if (find_status == mod2con.end()) { - return static_cast(HSA_EXT_STATUS_ERROR_INVALID_MODULE); - } - - HSAIL_ASM::BrigContainer *brig_container = - reinterpret_cast(find_status->second); - assert(brig_container); - delete brig_container; - mod2con.erase(find_status); - - return HSA_STATUS_SUCCESS; -} - -hsa_status_t ModuleValidate( - hsa_ext_module_t module, - uint32_t *result -) { - if (!result) { - return HSA_STATUS_ERROR_INVALID_ARGUMENT; - } - - auto find_status = mod2con.find(module); - if (find_status == mod2con.end()) { - return static_cast(HSA_EXT_STATUS_ERROR_INVALID_MODULE); - } - - HSAIL_ASM::BrigContainer *brig_container = - reinterpret_cast(find_status->second); - assert(brig_container); - HSAIL_ASM::Validator brig_validator(*brig_container); - *result = brig_validator.validate() ? 0 : 1; - - return HSA_STATUS_SUCCESS; -} - -hsa_status_t ModuleDisassemble( - hsa_ext_module_t module, - const char *hsail_text_filename -) { - if (!hsail_text_filename) { - return HSA_STATUS_ERROR_INVALID_ARGUMENT; - } - - auto find_status = mod2con.find(module); - if (find_status == mod2con.end()) { - return static_cast(HSA_EXT_STATUS_ERROR_INVALID_MODULE); - } - - HSAIL_ASM::BrigContainer *brig_container = - reinterpret_cast(find_status->second); - assert(brig_container); - HSAIL_ASM::Disassembler brig_disassembler(*brig_container); - if (brig_disassembler.run(hsail_text_filename)) { - return HSA_STATUS_ERROR_INVALID_ARGUMENT; - } - - return HSA_STATUS_SUCCESS; -} diff --git a/samples/common/assemble.hpp b/samples/common/assemble.hpp deleted file mode 100644 index 9024acca36..0000000000 --- a/samples/common/assemble.hpp +++ /dev/null @@ -1,50 +0,0 @@ -#ifndef ASSEMBLE_HPP_ -#define ASSEMBLE_HPP_ - -#include -#include "hsa.h" -#include "hsa_ext_finalize.h" - -#if defined(_MSC_VER) - #ifdef __cplusplus - extern "C" { - #endif -#endif - -hsa_status_t ModuleCreateFromHsailTextFile( - const char *hsail_text_filename, - hsa_ext_module_t *module -); - -hsa_status_t ModuleCreateFromBrigFile( - const char *hsail_text_filename, - hsa_ext_module_t *module -); - -hsa_status_t ModuleCreateFromHsailString( - const char *hsail_string, - hsa_ext_module_t *module -); - -hsa_status_t ModuleDestroy( - hsa_ext_module_t module -); - -hsa_status_t ModuleValidate( - hsa_ext_module_t module, - uint32_t *result -); - -hsa_status_t ModuleDisassemble( - hsa_ext_module_t module, - const char *hsail_text_filename -); - -#if defined(_MSC_VER) - #ifdef __cplusplus - } - #endif -#endif - - -#endif // ASSEMBLE_HPP_ diff --git a/samples/common/hsa_base_util.cpp b/samples/common/hsa_base_util.cpp index 4f3faa4338..a75fedcf98 100644 --- a/samples/common/hsa_base_util.cpp +++ b/samples/common/hsa_base_util.cpp @@ -55,8 +55,11 @@ bool HSA_UTIL::HsaInit() err = hsa_queue_create(device, queue_size, HSA_QUEUE_TYPE_MULTI, NULL, NULL, 0, 0, &command_queue); check(Creating the queue, err); - err = ModuleCreateFromHsailTextFile(hail_file_name, &module); - check(Module cration from hsail string, err); + if (!tool.assembleFromFile(hail_file_name)) { + std::cout << tool.output(); + return false; + } + module = tool.brigModule(); // Create hsail program. err = hsa_ext_program_create(HSA_MACHINE_MODEL_LARGE, HSA_PROFILE_FULL, HSA_DEFAULT_FLOAT_ROUNDING_MODE_ZERO, NULL, &hsa_program); diff --git a/samples/common/hsa_base_util.h b/samples/common/hsa_base_util.h index f9a0f213a3..7b1077c338 100644 --- a/samples/common/hsa_base_util.h +++ b/samples/common/hsa_base_util.h @@ -8,8 +8,8 @@ #include "hsa_ext_amd.h" #include "hsatimer.h" #include "utilities.h" -#include "assemble.hpp" #include "common.hpp" +#include "HSAILTool.h" @@ -35,6 +35,7 @@ class HSA_UTIL{ char hsa_kernel_name[128]; hsa_queue_t* command_queue; + HSAIL_ASM::Tool tool; hsa_ext_module_t module; hsa_ext_program_t hsa_program; hsa_executable_t hsaExecutable; diff --git a/samples/common/hsa_rsrc_factory.cpp b/samples/common/hsa_rsrc_factory.cpp index 34f47c4f78..3bbba59e8e 100755 --- a/samples/common/hsa_rsrc_factory.cpp +++ b/samples/common/hsa_rsrc_factory.cpp @@ -12,7 +12,6 @@ #include "hsa_rsrc_factory.hpp" #include "hsa_ext_finalize.h" -#include "assemble.hpp" #include "common.hpp" using namespace std; @@ -226,8 +225,11 @@ bool HsaRsrcFactory::LoadAndFinalize(AgentInfo *agent_info, // Copy handle of Brig object hsa_ext_module_t brig_module_v3; - status = ModuleCreateFromHsailTextFile(brig_path, &brig_module_v3); - check("Error in creating module from hsail text", status); + if (!tool.assembleFromFile(brig_path)) { + std::cout << tool.output(); + return false; + } + brig_module_v3 = tool.brigModule(); // Create hsail program. hsa_ext_program_t hsailProgram; diff --git a/samples/common/hsa_rsrc_factory.hpp b/samples/common/hsa_rsrc_factory.hpp index e685d4a4fc..88fc06c87d 100755 --- a/samples/common/hsa_rsrc_factory.hpp +++ b/samples/common/hsa_rsrc_factory.hpp @@ -13,6 +13,7 @@ #include "hsatimer.h" #include "hsa.h" #include "hsa_ext_finalize.h" +#include "HSAILTool.h" #define HSA_ARGUMENT_ALIGN_BYTES 16 @@ -270,6 +271,7 @@ class HsaRsrcFactory { // Maps an index for the user argument static uint32_t GetArgIndex(char *arg_value); + HSAIL_ASM::Tool tool; }; #endif // HSA_RSRC_FACTORY_H_