From b01abf4e3f6d9082a7c4a8fc34bb26f299b10425 Mon Sep 17 00:00:00 2001 From: "Ramesh Errabolu (xN/A) TX" Date: Fri, 13 Mar 2015 11:46:58 -0500 Subject: [PATCH] ECR #333755 - Update code to use 1.0F Api's of Finalizer. Presently disabling the build of all samples except BinarySearch [git-p4: depot-paths = "//depot/stg/hsa/drivers/hsa/runtime/": change = 1130579] [ROCm/ROCR-Runtime commit: f92e289a01760d588a8f00c4ac13459565d8b579] --- projects/rocr-runtime/samples/Makefile | 23 +- projects/rocr-runtime/samples/common/Makefile | 7 - .../{hsailUtil => common}/assemble.cpp | 40 +++- .../{hsailUtil => common}/assemble.hpp | 17 +- .../samples/{hsailUtil => common}/common.cpp | 0 .../samples/{hsailUtil => common}/common.hpp | 0 .../rocr-runtime/samples/common/elf_utils.cpp | 222 ------------------ .../rocr-runtime/samples/common/elf_utils.h | 41 ---- .../rocr-runtime/samples/common/hsa_base.cpp | 29 ++- .../rocr-runtime/samples/common/hsa_base.h | 8 +- .../samples/common/hsa_rsrc_factory.cpp | 68 ++---- .../samples/common/hsa_rsrc_factory.hpp | 4 +- .../rocr-runtime/samples/common/utilities.cpp | 4 +- .../rocr-runtime/samples/common/utilities.h | 4 +- .../rocr-runtime/samples/hsailUtil/Makefile | 7 - .../samples/hsailUtil/build/Makefile | 8 - .../samples/hsailUtil/build/Makefile.common | 51 ---- 17 files changed, 102 insertions(+), 431 deletions(-) delete mode 100644 projects/rocr-runtime/samples/common/Makefile rename projects/rocr-runtime/samples/{hsailUtil => common}/assemble.cpp (73%) rename projects/rocr-runtime/samples/{hsailUtil => common}/assemble.hpp (60%) rename projects/rocr-runtime/samples/{hsailUtil => common}/common.cpp (100%) rename projects/rocr-runtime/samples/{hsailUtil => common}/common.hpp (100%) delete mode 100644 projects/rocr-runtime/samples/common/elf_utils.cpp delete mode 100644 projects/rocr-runtime/samples/common/elf_utils.h delete mode 100644 projects/rocr-runtime/samples/hsailUtil/Makefile delete mode 100644 projects/rocr-runtime/samples/hsailUtil/build/Makefile delete mode 100644 projects/rocr-runtime/samples/hsailUtil/build/Makefile.common diff --git a/projects/rocr-runtime/samples/Makefile b/projects/rocr-runtime/samples/Makefile index 86ab8fdb1c..863345c050 100644 --- a/projects/rocr-runtime/samples/Makefile +++ b/projects/rocr-runtime/samples/Makefile @@ -2,17 +2,16 @@ OPENCL_DEPTH = ../.. include $(OPENCL_DEPTH)/hsadefs -SUBDIRS = NBody -SUBDIRS += DwtHarr1D -SUBDIRS += BitionicSort -SUBDIRS += BinarySearch -SUBDIRS += BinarySearch -SUBDIRS += BlackScholes -SUBDIRS += FloydWarshall -SUBDIRS += FastWalshTransform -SUBDIRS += MatrixTranspose -SUBDIRS += MatrixMultiplication -SUBDIRS += MonteCarloAsian -SUBDIRS += SimpleConvolution +#SUBDIRS = NBody +#SUBDIRS += DwtHarr1D +#SUBDIRS += BitionicSort +SUBDIRS = BinarySearch +#SUBDIRS += BlackScholes +#SUBDIRS += FloydWarshall +#SUBDIRS += FastWalshTransform +#SUBDIRS += MatrixTranspose +#SUBDIRS += MatrixMultiplication +#SUBDIRS += MonteCarloAsian +#SUBDIRS += SimpleConvolution include $(OPENCL_DEPTH)/hsarules diff --git a/projects/rocr-runtime/samples/common/Makefile b/projects/rocr-runtime/samples/common/Makefile deleted file mode 100644 index ae780dfdcf..0000000000 --- a/projects/rocr-runtime/samples/common/Makefile +++ /dev/null @@ -1,7 +0,0 @@ -OPENCL_DEPTH = ../../.. - -include $(OPENCL_DEPTH)/runtimenew/runtimedefs - -SUBDIRS = build - -include $(OPENCL_DEPTH)/runtimenew/runtimerules diff --git a/projects/rocr-runtime/samples/hsailUtil/assemble.cpp b/projects/rocr-runtime/samples/common/assemble.cpp similarity index 73% rename from projects/rocr-runtime/samples/hsailUtil/assemble.cpp rename to projects/rocr-runtime/samples/common/assemble.cpp index a79fb9e49d..f8b4ba8b8e 100644 --- a/projects/rocr-runtime/samples/hsailUtil/assemble.cpp +++ b/projects/rocr-runtime/samples/common/assemble.cpp @@ -8,11 +8,12 @@ #include #include "assemble.hpp" #include "hsa.h" -#include "hsa_ext_alt_finalize.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; @@ -20,7 +21,7 @@ namespace { hsa_status_t ModuleCreateFromHsailTextFile( const char *hsail_text_filename, - hsa_ext_alt_module_t *module + hsa_ext_module_t *module ) { if (!hsail_text_filename) { return HSA_STATUS_ERROR_INVALID_ARGUMENT; @@ -41,9 +42,28 @@ hsa_status_t ModuleCreateFromHsailTextFile( 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( + reinterpret_cast(brig_container->getBrigModule()), + reinterpret_cast(brig_container) + ) + ); + assert(insert_status.second); + module->handle = reinterpret_cast(brig_container->getBrigModule()); + return HSA_STATUS_SUCCESS; +} + hsa_status_t ModuleCreateFromHsailString( const char *hsail_string, - hsa_ext_alt_module_t *module + hsa_ext_module_t *module ) { if (!hsail_string || !module) { return HSA_STATUS_ERROR_INVALID_ARGUMENT; @@ -63,7 +83,7 @@ hsa_status_t ModuleCreateFromHsailString( hsail_parser.parseSource(); } catch (const SyntaxError) { delete brig_container; - return static_cast(HSA_EXT_ALT_STATUS_ERROR_INVALID_MODULE); + return static_cast(HSA_EXT_STATUS_ERROR_INVALID_MODULE); } try { @@ -84,11 +104,11 @@ hsa_status_t ModuleCreateFromHsailString( } hsa_status_t ModuleDestroy( - hsa_ext_alt_module_t module + hsa_ext_module_t module ) { auto find_status = mod2con.find(module.handle); if (find_status == mod2con.end()) { - return static_cast(HSA_EXT_ALT_STATUS_ERROR_INVALID_MODULE); + return static_cast(HSA_EXT_STATUS_ERROR_INVALID_MODULE); } HSAIL_ASM::BrigContainer *brig_container = @@ -101,7 +121,7 @@ hsa_status_t ModuleDestroy( } hsa_status_t ModuleValidate( - hsa_ext_alt_module_t module, + hsa_ext_module_t module, uint32_t *result ) { if (!result) { @@ -110,7 +130,7 @@ hsa_status_t ModuleValidate( auto find_status = mod2con.find(module.handle); if (find_status == mod2con.end()) { - return static_cast(HSA_EXT_ALT_STATUS_ERROR_INVALID_MODULE); + return static_cast(HSA_EXT_STATUS_ERROR_INVALID_MODULE); } HSAIL_ASM::BrigContainer *brig_container = @@ -123,7 +143,7 @@ hsa_status_t ModuleValidate( } hsa_status_t ModuleDisassemble( - hsa_ext_alt_module_t module, + hsa_ext_module_t module, const char *hsail_text_filename ) { if (!hsail_text_filename) { @@ -132,7 +152,7 @@ hsa_status_t ModuleDisassemble( auto find_status = mod2con.find(module.handle); if (find_status == mod2con.end()) { - return static_cast(HSA_EXT_ALT_STATUS_ERROR_INVALID_MODULE); + return static_cast(HSA_EXT_STATUS_ERROR_INVALID_MODULE); } HSAIL_ASM::BrigContainer *brig_container = diff --git a/projects/rocr-runtime/samples/hsailUtil/assemble.hpp b/projects/rocr-runtime/samples/common/assemble.hpp similarity index 60% rename from projects/rocr-runtime/samples/hsailUtil/assemble.hpp rename to projects/rocr-runtime/samples/common/assemble.hpp index c43f8aaac7..ed102c76a3 100644 --- a/projects/rocr-runtime/samples/hsailUtil/assemble.hpp +++ b/projects/rocr-runtime/samples/common/assemble.hpp @@ -3,29 +3,34 @@ #include #include "hsa.h" -#include "hsa_ext_alt_finalize.h" +#include "hsa_ext_finalize.h" hsa_status_t ModuleCreateFromHsailTextFile( const char *hsail_text_filename, - hsa_ext_alt_module_t *module + 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_alt_module_t *module + hsa_ext_module_t *module ); hsa_status_t ModuleDestroy( - hsa_ext_alt_module_t module + hsa_ext_module_t module ); hsa_status_t ModuleValidate( - hsa_ext_alt_module_t module, + hsa_ext_module_t module, uint32_t *result ); hsa_status_t ModuleDisassemble( - hsa_ext_alt_module_t module, + hsa_ext_module_t module, const char *hsail_text_filename ); diff --git a/projects/rocr-runtime/samples/hsailUtil/common.cpp b/projects/rocr-runtime/samples/common/common.cpp similarity index 100% rename from projects/rocr-runtime/samples/hsailUtil/common.cpp rename to projects/rocr-runtime/samples/common/common.cpp diff --git a/projects/rocr-runtime/samples/hsailUtil/common.hpp b/projects/rocr-runtime/samples/common/common.hpp similarity index 100% rename from projects/rocr-runtime/samples/hsailUtil/common.hpp rename to projects/rocr-runtime/samples/common/common.hpp diff --git a/projects/rocr-runtime/samples/common/elf_utils.cpp b/projects/rocr-runtime/samples/common/elf_utils.cpp deleted file mode 100644 index ae6a4aa2d9..0000000000 --- a/projects/rocr-runtime/samples/common/elf_utils.cpp +++ /dev/null @@ -1,222 +0,0 @@ -/* Copyright 2014 HSA Foundation Inc. All Rights Reserved. - * - * HSAF is granting you permission to use this software and documentation (if - * any) (collectively, the "Materials") pursuant to the terms and conditions - * of the Software License Agreement included with the Materials. If you do - * not have a copy of the Software License Agreement, contact the HSA Foundation for a copy. - * Redistribution and use in source and binary forms, with or without - * modification, are permitted provided that the following conditions - * are met: - * 1. Redistributions of source code must retain the above copyright - * notice, this list of conditions and the following disclaimer. - * 2. Redistributions in binary form must reproduce the above copyright - * notice, this list of conditions and the following disclaimer in the - * documentation and/or other materials provided with the distribution - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, FITNESS - * FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING - * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS WITH THE SOFTWARE. - */ - -#include -#include -#include -#include -#include "hsa.h" -#include "elf_utils.h" -#include "hsa_ext_finalize.h" - -enum { - SECTION_HSA_DATA = 0, - SECTION_HSA_CODE, - SECTION_HSA_OPERAND, -}; - -typedef struct SectionDesc SectionDesc; -struct SectionDesc { - int sectionId; - const char *brigName; - const char *bifName; -} - -sectionDescs[] = { - { SECTION_HSA_DATA, "hsa_data",".brig_hsa_data" }, - { SECTION_HSA_CODE, "hsa_code",".brig_hsa_code" }, - { SECTION_HSA_OPERAND,"hsa_operand",".brig_hsa_operand"}, -}; - -extern int fileno(FILE* stream); - -const SectionDesc* get_section_desc(int sectionId) { - const int NUM_PREDEFINED_SECTIONS = sizeof(sectionDescs)/sizeof(sectionDescs[0]); - for(int i=0; id_buf + shdr->sh_name; - if (sectionName && - ((strcmp(sectionName, desc->brigName) == 0) || - (strcmp(sectionName, desc->bifName) == 0))) { - return scn; - } - } - - return NULL; -} - -/* Extract section and copy into HsaBrig */ -static status_t extract_section_and_copy (Elf *elfP, - Elf_Data *secHdr, - const SectionDesc* desc, - hsa_ext_brig_module_t* brig_module, - hsa_ext_brig_section_id_t section_id) { - Elf_Scn* scn = NULL; - Elf_Data* data = NULL; - void* address_to_copy; - size_t section_size=0; - - scn = extract_elf_section(elfP, secHdr, desc); - - if (scn) { - if ((data = elf_getdata(scn, NULL)) == NULL) { - return STATUS_UNKNOWN; - } - section_size = data->d_size; - if (section_size > 0) { - address_to_copy = malloc(section_size); - memcpy(address_to_copy, data->d_buf, section_size); - } - } - - if ((!scn || section_size == 0)) { - return STATUS_UNKNOWN; - } - - /* Create a section header */ - brig_module->section[section_id] = (hsa_ext_brig_section_header_t*) address_to_copy; - - return STATUS_SUCCESS; -} - -/* Reads binary of BRIG and BIF format */ -status_t read_binary(hsa_ext_brig_module_t **brig_module_t, FILE* binary) { - /* Create the brig_module */ - uint32_t number_of_sections = 3; - hsa_ext_brig_module_t* brig_module; - - brig_module = (hsa_ext_brig_module_t*) - (malloc (sizeof(hsa_ext_brig_module_t) + sizeof(void*)*number_of_sections)); - brig_module->section_count = number_of_sections; - - status_t status; - Elf* elfP = NULL; - Elf32_Ehdr* ehdr = NULL; - Elf_Data *secHdr = NULL; - Elf_Scn* scn = NULL; - int fd; - - if (elf_version ( EV_CURRENT ) == EV_NONE) { - return STATUS_KERNEL_ELF_INITIALIZATION_FAILED; - } - - fd = fileno(binary); - if ((elfP = elf_begin(fd, ELF_C_READ, (Elf *)0)) == NULL) { - return STATUS_KERNEL_INVALID_ELF_CONTAINER; - } - - if (elf_kind (elfP) != ELF_K_ELF) { - return STATUS_KERNEL_INVALID_ELF_CONTAINER; - } - - if (((ehdr = elf32_getehdr(elfP)) == NULL) || - ((scn = elf_getscn(elfP, ehdr->e_shstrndx)) == NULL) || - ((secHdr = elf_getdata(scn, NULL)) == NULL)) { - return STATUS_KERNEL_INVALID_SECTION_HEADER; - } - - status = extract_section_and_copy(elfP, - secHdr, - get_section_desc(SECTION_HSA_DATA), - brig_module, - HSA_EXT_BRIG_SECTION_DATA); - - if (status != STATUS_SUCCESS) { - return STATUS_KERNEL_MISSING_DATA_SECTION; - } - - status = extract_section_and_copy(elfP, - secHdr, - get_section_desc(SECTION_HSA_CODE), - brig_module, - HSA_EXT_BRIG_SECTION_CODE); - - if (status != STATUS_SUCCESS) { - return STATUS_KERNEL_MISSING_CODE_SECTION; - } - - status = extract_section_and_copy(elfP, - secHdr, - get_section_desc(SECTION_HSA_OPERAND), - brig_module, - HSA_EXT_BRIG_SECTION_OPERAND); - - if (status != STATUS_SUCCESS) { - return STATUS_KERNEL_MISSING_OPERAND_SECTION; - } - - elf_end(elfP); - *brig_module_t = brig_module; - - return STATUS_SUCCESS; -} - -status_t create_brig_module_from_brig_file(const char* file_name, hsa_ext_brig_module_t** brig_module) { - FILE *fp = fopen(file_name, "rb"); - - status_t status = read_binary(brig_module, fp); - - if (status != STATUS_SUCCESS) { - printf("Could not create BRIG module: %d\n", status); - if (status == STATUS_KERNEL_INVALID_SECTION_HEADER || - status == STATUS_KERNEL_ELF_INITIALIZATION_FAILED || - status == STATUS_KERNEL_INVALID_ELF_CONTAINER) { - printf("The ELF file is invalid or possibley corrupted.\n"); - } - if (status == STATUS_KERNEL_MISSING_DATA_SECTION || - status == STATUS_KERNEL_MISSING_CODE_SECTION || - status == STATUS_KERNEL_MISSING_OPERAND_SECTION) { - printf("One or more ELF sections are missing. Use readelf command to \ - to check if hsa_data, hsa_code and hsa_operands exist.\n"); - } - } - - fclose(fp); - - return status; -} - -void destroy_brig_module(hsa_ext_brig_module_t* brig_module) { - for (int i=0; isection_count; i++) { - free (brig_module->section[i]); - } - free (brig_module); -} diff --git a/projects/rocr-runtime/samples/common/elf_utils.h b/projects/rocr-runtime/samples/common/elf_utils.h deleted file mode 100644 index 3305d33c00..0000000000 --- a/projects/rocr-runtime/samples/common/elf_utils.h +++ /dev/null @@ -1,41 +0,0 @@ -/* Copyright 2014 HSA Foundation Inc. All Rights Reserved. - * - * HSAF is granting you permission to use this software and documentation (if - * any) (collectively, the "Materials") pursuant to the terms and conditions - * of the Software License Agreement included with the Materials. If you do - * not have a copy of the Software License Agreement, contact the HSA Foundation for a copy. - * Redistribution and use in source and binary forms, with or without - * modification, are permitted provided that the following conditions - * are met: - * 1. Redistributions of source code must retain the above copyright - * notice, this list of conditions and the following disclaimer. - * 2. Redistributions in binary form must reproduce the above copyright - * notice, this list of conditions and the following disclaimer in the - * documentation and/or other materials provided with the distribution - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, FITNESS - * FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING - * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS WITH THE SOFTWARE. - */ - -#pragma once - -#include "hsa_ext_finalize.h" - -//typedef enum status_t status_t; -typedef enum status_t { - STATUS_SUCCESS=0, - STATUS_KERNEL_INVALID_SECTION_HEADER=1, - STATUS_KERNEL_ELF_INITIALIZATION_FAILED=2, - STATUS_KERNEL_INVALID_ELF_CONTAINER=3, - STATUS_KERNEL_MISSING_DATA_SECTION=4, - STATUS_KERNEL_MISSING_CODE_SECTION=5, - STATUS_KERNEL_MISSING_OPERAND_SECTION=6, - STATUS_UNKNOWN=7, -} status_t; - -status_t create_brig_module_from_brig_file(const char* file_name, hsa_ext_brig_module_t** brig_module); - -void destroy_brig_module(hsa_ext_brig_module_t* brig_module); diff --git a/projects/rocr-runtime/samples/common/hsa_base.cpp b/projects/rocr-runtime/samples/common/hsa_base.cpp index e451023c5e..30664decc9 100644 --- a/projects/rocr-runtime/samples/common/hsa_base.cpp +++ b/projects/rocr-runtime/samples/common/hsa_base.cpp @@ -1,10 +1,12 @@ #include "hsa_base.h" +#if 0 void HSA::SetBrigFileAndKernelName(char * brig_file_name, char *kernel_name) { strcpy(hsa_brig_file_name, brig_file_name); strcpy(hsa_kernel_name, kernel_name); } +#endif HSA::HSA() { @@ -16,7 +18,7 @@ HSA::~HSA() } - +#if 0 bool HSA::HsaInit() { err = hsa_init(); @@ -62,29 +64,29 @@ double HSA::Run(int dim, int group_x, int group_y, int group_z, int s_size, int check(Creating the brig module from the input brig file, err); // Copy handle of Brig object - hsa_ext_alt_module_t brig_module_v3; + hsa_ext_module_t brig_module_v3; brig_module_v3.handle = uint64_t(local_brig_module); // Create hsail program. - hsa_ext_alt_program_t local_hsa_program; - err = hsa_ext_alt_program_create(HSA_MACHINE_MODEL_LARGE, + hsa_ext_program_t local_hsa_program; + err = hsa_ext_program_create(HSA_MACHINE_MODEL_LARGE, HSA_PROFILE_FULL, HSA_DEFAULT_FLOAT_ROUNDING_MODE_ZERO, NULL, &local_hsa_program); check("Error in creating program object", err); // Add hsail module. - err = hsa_ext_alt_program_add_module(local_hsa_program, brig_module_v3); + err = hsa_ext_program_add_module(local_hsa_program, brig_module_v3); check("Error in adding module to program object", err); // Finalize hsail program. hsa_isa_t isa; memset(&isa, 0, sizeof(hsa_isa_t)); - hsa_ext_alt_control_directives_t control_directives; - memset(&control_directives, 0, sizeof(hsa_ext_alt_control_directives_t)); + hsa_ext_control_directives_t control_directives; + memset(&control_directives, 0, sizeof(hsa_ext_control_directives_t)); hsa_code_object_t code_object; - err = hsa_ext_alt_program_finalize(local_hsa_program, + err = hsa_ext_program_finalize(local_hsa_program, isa, 0, control_directives, @@ -93,7 +95,7 @@ double HSA::Run(int dim, int group_x, int group_y, int group_z, int s_size, int &code_object); check("Error in finalizing program object", err); - //status = hsa_ext_alt_program_destroy(hsailProgram); + //status = hsa_ext_program_destroy(hsailProgram); //check("Error in destroying program object", status); // Create executable. @@ -158,13 +160,14 @@ double HSA::Run(int dim, int group_x, int group_y, int group_z, int s_size, int check(Finding a kernarg memory region, err); void* local_kernel_arg_buffer = NULL; - size_t local_kernel_arg_buffer_size; - hsa_executable_symbol_get_info(kernelSymbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE, &local_kernel_arg_buffer_size); + //size_t local_kernel_arg_buffer_size; + //hsa_executable_symbol_get_info(kernelSymbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE, &local_kernel_arg_buffer_size); /* * Allocate the kernel argument buffer from the correct region. */ - err = hsa_memory_allocate(local_kernarg_region, local_kernel_arg_buffer_size, &local_kernel_arg_buffer); + //err = hsa_memory_allocate(local_kernarg_region, local_kernel_arg_buffer_size, &local_kernel_arg_buffer); + err = hsa_memory_allocate(local_kernarg_region, kernel_args_size, &local_kernel_arg_buffer); check(Allocating kernel argument memory buffer, err); memcpy(local_kernel_arg_buffer, kernel_args, kernel_args_size); local_dispatch_packet.kernarg_address = local_kernel_arg_buffer; @@ -238,3 +241,5 @@ void HSA::Close() check(Shutting down the runtime, err); } +#endif + diff --git a/projects/rocr-runtime/samples/common/hsa_base.h b/projects/rocr-runtime/samples/common/hsa_base.h index 356e4857fe..e1f66e284d 100644 --- a/projects/rocr-runtime/samples/common/hsa_base.h +++ b/projects/rocr-runtime/samples/common/hsa_base.h @@ -5,8 +5,6 @@ #include #include "hsa.h" #include "hsa_ext_finalize.h" -#include "hsa_ext_alt_finalize.h" -#include "elf_utils.h" #include "hsatimer.h" #include "utilities.h" @@ -14,14 +12,15 @@ class HSA{ public: HSA(); ~HSA(); - +#if 0 public: void SetBrigFileAndKernelName(char *brig_file_name, char *kernel_name); bool HsaInit(); void Close(); double Run(int dim, int group_x, int group_y, int group_z, int s_size, int grid_x, int grid_y, int grid_z, void* kernel_args, int kernel_args_size); - +#endif public: + #if 0 hsa_status_t err; uint32_t queue_size; hsa_agent_t device; @@ -37,6 +36,7 @@ class HSA{ hsa_ext_code_descriptor_t *hsa_code_descriptor; hsa_kernel_dispatch_packet_t dispatch_packet; // needs to be set manually each time hsa_region_t hsa_kernarg_region; + #endif }; diff --git a/projects/rocr-runtime/samples/common/hsa_rsrc_factory.cpp b/projects/rocr-runtime/samples/common/hsa_rsrc_factory.cpp index 2937100690..c54b40774f 100755 --- a/projects/rocr-runtime/samples/common/hsa_rsrc_factory.cpp +++ b/projects/rocr-runtime/samples/common/hsa_rsrc_factory.cpp @@ -10,9 +10,11 @@ #include #include "hsa.h" -#include "elf_utils.h" #include "hsa_rsrc_factory.hpp" -#include "hsa_ext_alt_finalize.h" +#include "hsa_ext_finalize.h" + +#include "assemble.hpp" +#include "common.hpp" using namespace std; @@ -71,40 +73,6 @@ static hsa_status_t get_gpu_agents(hsa_agent_t agent, void *data) { return HSA_STATUS_SUCCESS; } -// Finds the specified symbols offset in the specified brig_module. -// If the symbol is found the function returns HSA_STATUS_SUCCESS, -// otherwise it returns HSA_STATUS_ERROR. -hsa_status_t hsa_find_symbol_offset(hsa_ext_brig_module_t *brig_module, - char *symbol_name, - hsa_ext_brig_code_section_offset32_t *offset) { - - // Get the data section - hsa_ext_brig_section_header_t *data_hdr = brig_module->section[HSA_EXT_BRIG_SECTION_DATA]; - - // Get the code section - hsa_ext_brig_section_header_t* code_hdr = brig_module->section[HSA_EXT_BRIG_SECTION_CODE]; - - // First entry into the BRIG code section - BrigCodeOffset32_t code_offset = code_hdr->header_byte_count; - BrigBase* code_entry = (BrigBase*) ((char*)code_hdr + code_offset); - while (code_offset != code_hdr->byte_count) { - if (code_entry->kind == BRIG_KIND_DIRECTIVE_KERNEL) { - - // Now find the data in the data section - BrigDirectiveExecutable* directive_kernel = (BrigDirectiveExecutable*) (code_entry); - BrigDataOffsetString32_t data_name_offset = directive_kernel->name; - BrigData* data_entry = (BrigData*)((char*) data_hdr + data_name_offset); - if (!strncmp(symbol_name, (char*) data_entry->bytes, strlen(symbol_name))) { - *offset = code_offset; - return HSA_STATUS_SUCCESS; - } - } - code_offset += code_entry->byteCount; - code_entry = (BrigBase*) ((char*)code_hdr + code_offset); - } - return HSA_STATUS_ERROR; -} - // Definitions for Static Data members of the class char* HsaRsrcFactory::brig_path_ = NULL; uint32_t HsaRsrcFactory::num_cus_; @@ -235,7 +203,9 @@ bool HsaRsrcFactory::LoadAndFinalize(AgentInfo *agent_info, const char *brig_path, char *kernel_name, hsa_executable_symbol_t *code_desc) { + hsa_status_t status; // Load BRIG, encapsulated in an ELF container, into a BRIG module. + /* status_t build_err; hsa_ext_brig_module_t *brig_obj; build_err = (status_t)create_brig_module_from_brig_file(brig_path, &brig_obj); @@ -246,32 +216,36 @@ bool HsaRsrcFactory::LoadAndFinalize(AgentInfo *agent_info, hsa_ext_brig_code_section_offset32_t kernel_symbol; status = hsa_find_symbol_offset(brig_obj, kernel_name, &kernel_symbol); check("Error in Finding the Symbol Offset for the Kernel", status); + */ // Copy handle of Brig object - hsa_ext_alt_module_t brig_module_v3; - brig_module_v3.handle = uint64_t(brig_obj); + hsa_ext_module_t brig_module_v3; + status = ModuleCreateFromHsailTextFile(brig_path, &brig_module_v3); + check("Error in creating module from hsail text", status); // Create hsail program. - hsa_ext_alt_program_t hsailProgram; - status = hsa_ext_alt_program_create(HSA_MACHINE_MODEL_LARGE, + hsa_ext_program_t hsailProgram; + status = hsa_ext_program_create(HSA_MACHINE_MODEL_LARGE, HSA_PROFILE_FULL, HSA_DEFAULT_FLOAT_ROUNDING_MODE_ZERO, NULL, &hsailProgram); check("Error in creating program object", status); // Add hsail module. - status = hsa_ext_alt_program_add_module(hsailProgram, brig_module_v3); + status = hsa_ext_program_add_module(hsailProgram, brig_module_v3); check("Error in adding module to program object", status); // Finalize hsail program. - hsa_isa_t isa; - memset(&isa, 0, sizeof(hsa_isa_t)); + hsa_isa_t isa = {0}; + status = hsa_agent_get_info(agent_info->dev_id, HSA_AGENT_INFO_ISA, &isa); + std::cout << "Value of device Isa Id: " << isa.handle << std::endl; + check("Error in getting Id of Isa supported by agent", status); - hsa_ext_alt_control_directives_t control_directives; - memset(&control_directives, 0, sizeof(hsa_ext_alt_control_directives_t)); + hsa_ext_control_directives_t control_directives; + memset(&control_directives, 0, sizeof(hsa_ext_control_directives_t)); hsa_code_object_t code_object; - status = hsa_ext_alt_program_finalize(hsailProgram, + status = hsa_ext_program_finalize(hsailProgram, isa, 0, control_directives, @@ -280,7 +254,7 @@ bool HsaRsrcFactory::LoadAndFinalize(AgentInfo *agent_info, &code_object); check("Error in finalizing program object", status); - //status = hsa_ext_alt_program_destroy(hsailProgram); + //status = hsa_ext_program_destroy(hsailProgram); //check("Error in destroying program object", status); // Create executable. diff --git a/projects/rocr-runtime/samples/common/hsa_rsrc_factory.hpp b/projects/rocr-runtime/samples/common/hsa_rsrc_factory.hpp index 2c0dce83eb..6764f020ab 100755 --- a/projects/rocr-runtime/samples/common/hsa_rsrc_factory.hpp +++ b/projects/rocr-runtime/samples/common/hsa_rsrc_factory.hpp @@ -22,7 +22,9 @@ #define check(msg, status) \ if (status != HSA_STATUS_SUCCESS) { \ - printf("%s\n", msg); \ + const char *emsg = 0; \ + hsa_status_string(status, &emsg); \ + printf("%s: %s\n", msg, emsg ? emsg : ""); \ exit(1); \ } diff --git a/projects/rocr-runtime/samples/common/utilities.cpp b/projects/rocr-runtime/samples/common/utilities.cpp index 96a5d34cda..82834bde90 100644 --- a/projects/rocr-runtime/samples/common/utilities.cpp +++ b/projects/rocr-runtime/samples/common/utilities.cpp @@ -81,7 +81,7 @@ int FillRandom( return 0; } - +#if 0 //get a memory region that can be used for global memory allocations. hsa_status_t get_global_region(hsa_region_t region, void* data) { @@ -143,6 +143,7 @@ hsa_status_t find_symbol_offset(hsa_ext_brig_module_t* brig_module, } return HSA_STATUS_ERROR; } +#endif /* * Determines if the given agent is of type HSA_DEVICE_TYPE_GPU @@ -167,6 +168,7 @@ hsa_status_t find_gpu(hsa_agent_t agent, void *data) return HSA_STATUS_SUCCESS; } + /* * Determines if a memory region can be used for kernarg * allocations. diff --git a/projects/rocr-runtime/samples/common/utilities.h b/projects/rocr-runtime/samples/common/utilities.h index 9239712caf..3c781bb780 100644 --- a/projects/rocr-runtime/samples/common/utilities.h +++ b/projects/rocr-runtime/samples/common/utilities.h @@ -6,7 +6,6 @@ #include "hsa.h" #include "hsa_ext_finalize.h" -#include "elf_utils.h" #include #include @@ -190,7 +189,8 @@ hsa_status_t get_global_region(hsa_region_t region, void* data); * If the symbol is found the function returns HSA_STATUS_SUCCESS, * otherwise it returns HSA_STATUS_ERROR. */ -hsa_status_t find_symbol_offset(hsa_ext_brig_module_t* brig_module, char* symbol_name, hsa_ext_brig_code_section_offset32_t* offset); + +//hsa_status_t find_symbol_offset(hsa_ext_brig_module_t* brig_module, char* symbol_name, hsa_ext_brig_code_section_offset32_t* offset); /* * Determines if the given agent is of type HSA_DEVICE_TYPE_GPU diff --git a/projects/rocr-runtime/samples/hsailUtil/Makefile b/projects/rocr-runtime/samples/hsailUtil/Makefile deleted file mode 100644 index ae780dfdcf..0000000000 --- a/projects/rocr-runtime/samples/hsailUtil/Makefile +++ /dev/null @@ -1,7 +0,0 @@ -OPENCL_DEPTH = ../../.. - -include $(OPENCL_DEPTH)/runtimenew/runtimedefs - -SUBDIRS = build - -include $(OPENCL_DEPTH)/runtimenew/runtimerules diff --git a/projects/rocr-runtime/samples/hsailUtil/build/Makefile b/projects/rocr-runtime/samples/hsailUtil/build/Makefile deleted file mode 100644 index 7e15928ce4..0000000000 --- a/projects/rocr-runtime/samples/hsailUtil/build/Makefile +++ /dev/null @@ -1,8 +0,0 @@ -OPENCL_DEPTH = ../../../.. - -include $(OPENCL_DEPTH)/runtimenew/runtimedefs - -BUILD_SUBDIRS = $(DEFAULT_TARGETS) -BUILD_MAKEFILE = Makefile.common - -include $(OPENCL_DEPTH)/runtimenew/runtimerules diff --git a/projects/rocr-runtime/samples/hsailUtil/build/Makefile.common b/projects/rocr-runtime/samples/hsailUtil/build/Makefile.common deleted file mode 100644 index 49695b9b8e..0000000000 --- a/projects/rocr-runtime/samples/hsailUtil/build/Makefile.common +++ /dev/null @@ -1,51 +0,0 @@ -include $(OPENCL_DEPTH)/hsadefs - -LIB_TARGET = testcommon -vpath %.cpp $(COMPONENT_DEPTH) -CPPFILES := $(notdir $(wildcard $(COMPONENT_DEPTH)/*.cpp)) - -ifdef ATI_BITS_64 - LIB_SUFFIX = 64 - NBITS = 64 -else - LIB_SUFFIX = - ifndef ATI_OS_WINDOWS - NBITS := 32 - endif -endif - -ifdef ATI_OS_WINDOWS - CORE_LIB = dll - LFLAGS += /subsystem:console - LIB_PREFIX = -else - CORE_LIB = so - LIB_PREFIX = lib -endif - -ifdef ATI_OS_LINUX - GCXXOPTS := $(filter-out -fno-rtti,$(GCXXOPTS)) - GCXXOPTS := $(filter-out -fno-exceptions,$(GCXXOPTS)) - LFLAGS += -L$(DIST_LIB_DEST) -lpthread $(LIBSTDCXX) -lm -ldl -lrt -endif - -export BUILD_HSA_TARGET=yes - -LCINCS := $(INCSWITCH) "$(OPENCL_DEPTH)/compiler/finalizer/HSAIL/hsail-tools/libHSAIL" -LCINCS += $(INCSWITCH) "$(OPENCL_DEPTH)/compiler/finalizer/HSAIL/hsail-tools/libHSAIL/$(FULL_BUILD_DIR)" -LCINCS += $(INCSWITCH) "$(OPENCL_DEPTH)/compiler/finalizer/Interface" -LCINCS += $(INCSWITCH) "$(OPENCL_DEPTH)/runtime/inc" - -LLLIBS := $(OPENCL_DEPTH)/compiler/finalizer/HSAIL/hsail-tools/libHSAIL/$(FULL_BUILD_DIR)/libhsail$(LIB_EXT) -LLLIBS += $(OPENCL_DEPTH)/contrib/gtest-1.6.0/$(FULL_BUILD_DIR)/libgtest$(LIB_EXT) -LLLIBS += $(OPENCL_DEPTH)/runtime/test/gcommon/$(FULL_BUILD_DIR)/gtestcommon$(LIB_EXT) - -RUNTIME_BUILD = build/$(OS_TYPE)/$(CORE_LIB)/$(BUILD_DIR) - -ifdef ATI_OS_LINUX - LFLAGS += -L$(OPENCL_DEPTH)/runtime/core/$(RUNTIME_BUILD) -lhsa-runtime$(LIB_SUFFIX) -else - LLLIBS += $(OPENCL_DEPTH)/runtime/core/$(RUNTIME_BUILD)/hsa-runtime$(LIB_SUFFIX)$(LIB_EXT) -endif - -include $(OPENCL_DEPTH)/hsarules