AI/rocm-systems
2
0
Forkuj 0
Kod Zgłoszenia Oczekujące zmiany Actions Packages Projekty Wydania Wiki Aktywność

ECR #333755 - Update code to use 1.0F Api's of Finalizer. Presently disabling the build of all samples except BinarySearch

Przeglądaj źródła 
[git-p4: depot-paths = "//depot/stg/hsa/drivers/hsa/runtime/": change = 1130579]


[ROCm/ROCR-Runtime commit: f92e289a01]
This commit is contained in:
Ramesh Errabolu (xN/A) TX
2015-03-13 11:46:58 -05:00
rodzic 3e8dfb3ac5
commit b01abf4e3f
17 zmienionych plików z 102 dodań i 431 usunięć
Ściągnij plik aktualizacji Ściągnij plik porównania Expand all files Collapse all files
projects/rocr-runtime/samples/Makefile
+11 -12
Wyświetl plik
@@ -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
projects/rocr-runtime/samples/common/Makefile
-7
Wyświetl plik
@@ -1,7 +0,0 @@
OPENCL_DEPTH = ../../..
include $(OPENCL_DEPTH)/runtimenew/runtimedefs
SUBDIRS = build
include $(OPENCL_DEPTH)/runtimenew/runtimerules
projects/rocr-runtime/samples/hsailUtil/assemble.cpp → projects/rocr-runtime/samples/common/assemble.cpp
+30 -10
Wyświetl plik
@@ -8,11 +8,12 @@
#include <utility>
#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<uint64_t, uint64_t> 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_status_t>(HSA_EXT_STATUS_ERROR_INVALID_MODULE); }
auto insert_status = mod2con.insert(
std::make_pair<uint64_t, uint64_t>(
reinterpret_cast<uint64_t>(brig_container->getBrigModule()),
reinterpret_cast<uint64_t>(brig_container)
)
);
assert(insert_status.second);
module->handle = reinterpret_cast<uint64_t>(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_status_t>(HSA_EXT_ALT_STATUS_ERROR_INVALID_MODULE);
return static_cast<hsa_status_t>(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_status_t>(HSA_EXT_ALT_STATUS_ERROR_INVALID_MODULE);
return static_cast<hsa_status_t>(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_status_t>(HSA_EXT_ALT_STATUS_ERROR_INVALID_MODULE);
return static_cast<hsa_status_t>(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_status_t>(HSA_EXT_ALT_STATUS_ERROR_INVALID_MODULE);
return static_cast<hsa_status_t>(HSA_EXT_STATUS_ERROR_INVALID_MODULE);
}
HSAIL_ASM::BrigContainer *brig_container =
projects/rocr-runtime/samples/hsailUtil/assemble.hpp → projects/rocr-runtime/samples/common/assemble.hpp
+11 -6
Wyświetl plik
@@ -3,29 +3,34 @@
#include <cstdint>
#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
);
projects/rocr-runtime/samples/hsailUtil/common.cpp → projects/rocr-runtime/samples/common/common.cpp
Wyświetl plik
projects/rocr-runtime/samples/hsailUtil/common.hpp → projects/rocr-runtime/samples/common/common.hpp
Wyświetl plik
projects/rocr-runtime/samples/common/elf_utils.cpp
-222
Wyświetl plik
@@ -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 <stdlib.h>
#include <libelf.h>
#include <string.h>
#include <stdio.h>
#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; i<NUM_PREDEFINED_SECTIONS; ++i) {
if (sectionDescs[i].sectionId == sectionId) {
return &sectionDescs[i];
}
}
return NULL;
}
static Elf_Scn* extract_elf_section (Elf *elfP,
Elf_Data *secHdr,
const SectionDesc* desc) {
int cnt = 0;
Elf_Scn* scn = NULL;
Elf32_Shdr* shdr = NULL;
char* sectionName = NULL;
/* Iterate thru the elf sections */
for (cnt = 1, scn = NULL; scn = elf_nextscn(elfP, scn); cnt++) {
if (((shdr = elf32_getshdr(scn)) == NULL)) {
return NULL;
}
sectionName = (char *)secHdr->d_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; i<brig_module->section_count; i++) {
free (brig_module->section[i]);
}
free (brig_module);
}
projects/rocr-runtime/samples/common/elf_utils.h
-41
Wyświetl plik
@@ -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);
projects/rocr-runtime/samples/common/hsa_base.cpp
+17 -12
Wyświetl plik
@@ -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
projects/rocr-runtime/samples/common/hsa_base.h
+4 -4
Wyświetl plik
@@ -5,8 +5,6 @@
#include <vector>
#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
};
projects/rocr-runtime/samples/common/hsa_rsrc_factory.cpp
+21 -47
Wyświetl plik
@@ -10,9 +10,11 @@
#include <string>
#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.
projects/rocr-runtime/samples/common/hsa_rsrc_factory.hpp
+3 -1
Wyświetl plik
@@ -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 : "<unknown error>"); \
exit(1); \
}
projects/rocr-runtime/samples/common/utilities.cpp
+3 -1
Wyświetl plik
@@ -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.
projects/rocr-runtime/samples/common/utilities.h
+2 -2
Wyświetl plik
@@ -6,7 +6,6 @@
#include "hsa.h"
#include "hsa_ext_finalize.h"
#include "elf_utils.h"
#include <string.h>
#include<iostream>
@@ -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
projects/rocr-runtime/samples/hsailUtil/Makefile
-7
Wyświetl plik
@@ -1,7 +0,0 @@
OPENCL_DEPTH = ../../..
include $(OPENCL_DEPTH)/runtimenew/runtimedefs
SUBDIRS = build
include $(OPENCL_DEPTH)/runtimenew/runtimerules
projects/rocr-runtime/samples/hsailUtil/build/Makefile
-8
Wyświetl plik
@@ -1,8 +0,0 @@
OPENCL_DEPTH = ../../../..
include $(OPENCL_DEPTH)/runtimenew/runtimedefs
BUILD_SUBDIRS = $(DEFAULT_TARGETS)
BUILD_MAKEFILE = Makefile.common
include $(OPENCL_DEPTH)/runtimenew/runtimerules
projects/rocr-runtime/samples/hsailUtil/build/Makefile.common
-51
Wyświetl plik
@@ -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