AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

ECR #333755 - Merging HSA samples from prm to stg to conform the 1.0F

Browse Source 
[git-p4: depot-paths = "//depot/stg/hsa/drivers/hsa/runtime/": change = 1138673]


[ROCm/ROCR-Runtime commit: 73a921b926]
This commit is contained in:
Ding, Wei (xN/A) TX
2015-04-08 14:46:33 -05:00
parent 1018e31ef7
commit c3f014c75d
3 changed files with 25 additions and 19 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rocr-runtime/samples/common/assemble.cpp
+10 -10
View File
@@ -16,7 +16,7 @@
#include "HSAILBrigObjectFile.h"
namespace {
std::unordered_map<uint64_t, uint64_t> mod2con;
std::unordered_map<BrigModule_t, uint64_t> mod2con;
} // namespace anonymous
hsa_status_t ModuleCreateFromHsailTextFile(
@@ -51,13 +51,13 @@ hsa_status_t ModuleCreateFromBrigFile(
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()),
std::make_pair<BrigModule_t, uint64_t>(
brig_container->getBrigModule(),
reinterpret_cast<uint64_t>(brig_container)
)
);
assert(insert_status.second);
module->handle = reinterpret_cast<uint64_t>(brig_container->getBrigModule());
*module = brig_container->getBrigModule();
return HSA_STATUS_SUCCESS;
}
@@ -88,8 +88,8 @@ hsa_status_t ModuleCreateFromHsailString(
try {
auto insert_status = mod2con.insert(
std::make_pair<uint64_t, uint64_t>(
reinterpret_cast<uint64_t>(brig_container->getBrigModule()),
std::make_pair<BrigModule_t, uint64_t>(
brig_container->getBrigModule(),
reinterpret_cast<uint64_t>(brig_container)
)
);
@@ -99,14 +99,14 @@ hsa_status_t ModuleCreateFromHsailString(
return HSA_STATUS_ERROR_OUT_OF_RESOURCES;
}
module->handle = reinterpret_cast<uint64_t>(brig_container->getBrigModule());
*module = brig_container->getBrigModule();
return HSA_STATUS_SUCCESS;
}
hsa_status_t ModuleDestroy(
hsa_ext_module_t module
) {
auto find_status = mod2con.find(module.handle);
auto find_status = mod2con.find(module);
if (find_status == mod2con.end()) {
return static_cast<hsa_status_t>(HSA_EXT_STATUS_ERROR_INVALID_MODULE);
}
@@ -128,7 +128,7 @@ hsa_status_t ModuleValidate(
return HSA_STATUS_ERROR_INVALID_ARGUMENT;
}
auto find_status = mod2con.find(module.handle);
auto find_status = mod2con.find(module);
if (find_status == mod2con.end()) {
return static_cast<hsa_status_t>(HSA_EXT_STATUS_ERROR_INVALID_MODULE);
}
@@ -150,7 +150,7 @@ hsa_status_t ModuleDisassemble(
return HSA_STATUS_ERROR_INVALID_ARGUMENT;
}
auto find_status = mod2con.find(module.handle);
auto find_status = mod2con.find(module);
if (find_status == mod2con.end()) {
return static_cast<hsa_status_t>(HSA_EXT_STATUS_ERROR_INVALID_MODULE);
}
projects/rocr-runtime/samples/common/hsa_base_util.cpp
+4 -5
View File
@@ -175,7 +175,6 @@ double HSA_UTIL::Run(int dim, int group_x, int group_y, int group_z, int s_size,
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);
@@ -191,8 +190,8 @@ double HSA_UTIL::Run(int dim, int group_x, int group_y, int group_z, int s_size,
/*
* Register the memory region for the argument buffer.
*/
//err = hsa_memory_register(kernel_args, kernel_args_size); ////////////////????????????????????
//check(Registering the argument buffer, err); /////////////???????????????
//err = hsa_memory_register(kernel_args, kernel_args_size);
//check(Registering the argument buffer, err);
/*
* Obtain the current queue write index.
@@ -229,8 +228,8 @@ double HSA_UTIL::Run(int dim, int group_x, int group_y, int group_z, int s_size,
* Cleanup all allocated resources.
*/
err = hsa_memory_free(local_kernel_arg_buffer);
check(Deallocate memory, err);
err = hsa_memory_free(local_kernel_arg_buffer);
check(Deallocate memory, err);
err=hsa_signal_destroy(local_signal);
check(Destroying the local_signal, err);
projects/rocr-runtime/samples/common/hsa_base_util.h
+11 -4
View File
@@ -1,5 +1,6 @@
#ifndef __HSA_UTIL__
#define __HSA_UTIL__
#ifndef __HSA_BASE__
#define __HSA_BASE__
#include <vector>
#include "hsa.h"
@@ -10,6 +11,8 @@
#include "assemble.hpp"
#include "common.hpp"
class HSA_UTIL{
public:
HSA_UTIL();
@@ -18,8 +21,8 @@ class HSA_UTIL{
public:
void GetHsailNameAndKernelName(char *hail_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);
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);
public:
hsa_status_t err;
@@ -31,6 +34,10 @@ class HSA_UTIL{
hsa_queue_t* command_queue;
hsa_signal_t hsa_signal;
//hsa_ext_brig_module_t* brig_module;
//hsa_ext_brig_module_handle_t module;
//hsa_ext_program_handle_t hsa_program;
//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;
};