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]
Tá an tiomantas seo le fáil i:
@@ -1,167 +0,0 @@
|
||||
#include <cassert>
|
||||
#include <cstdint>
|
||||
#include <fstream>
|
||||
#include <new>
|
||||
#include <sstream>
|
||||
#include <string>
|
||||
#include <unordered_map>
|
||||
#include <utility>
|
||||
#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<BrigModule_t, uint64_t> 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_status_t>(HSA_EXT_STATUS_ERROR_INVALID_MODULE); }
|
||||
auto insert_status = mod2con.insert(
|
||||
std::make_pair<BrigModule_t, uint64_t>(
|
||||
brig_container->getBrigModule(),
|
||||
reinterpret_cast<uint64_t>(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_status_t>(HSA_EXT_STATUS_ERROR_INVALID_MODULE);
|
||||
}
|
||||
|
||||
try {
|
||||
auto insert_status = mod2con.insert(
|
||||
std::make_pair<BrigModule_t, uint64_t>(
|
||||
brig_container->getBrigModule(),
|
||||
reinterpret_cast<uint64_t>(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_status_t>(HSA_EXT_STATUS_ERROR_INVALID_MODULE);
|
||||
}
|
||||
|
||||
HSAIL_ASM::BrigContainer *brig_container =
|
||||
reinterpret_cast<HSAIL_ASM::BrigContainer*>(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_status_t>(HSA_EXT_STATUS_ERROR_INVALID_MODULE);
|
||||
}
|
||||
|
||||
HSAIL_ASM::BrigContainer *brig_container =
|
||||
reinterpret_cast<HSAIL_ASM::BrigContainer*>(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_status_t>(HSA_EXT_STATUS_ERROR_INVALID_MODULE);
|
||||
}
|
||||
|
||||
HSAIL_ASM::BrigContainer *brig_container =
|
||||
reinterpret_cast<HSAIL_ASM::BrigContainer*>(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;
|
||||
}
|
||||
@@ -1,50 +0,0 @@
|
||||
#ifndef ASSEMBLE_HPP_
|
||||
#define ASSEMBLE_HPP_
|
||||
|
||||
#include <cstdint>
|
||||
#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_
|
||||
@@ -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);
|
||||
|
||||
@@ -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;
|
||||
|
||||
@@ -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;
|
||||
|
||||
@@ -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_
|
||||
|
||||
Tagairt in Eagrán Nua
Cuir bac ar úsáideoir