From c3f014c75d2288806f238d8c92b359e66fcd6a18 Mon Sep 17 00:00:00 2001 From: "Ding, Wei (xN/A) TX" Date: Wed, 8 Apr 2015 14:46:33 -0500 Subject: [PATCH] ECR #333755 - Merging HSA samples from prm to stg to conform the 1.0F [git-p4: depot-paths = "//depot/stg/hsa/drivers/hsa/runtime/": change = 1138673] [ROCm/ROCR-Runtime commit: 73a921b926e2c147c09646e2189a15749c1e067e] --- .../rocr-runtime/samples/common/assemble.cpp | 20 +++++++++---------- .../samples/common/hsa_base_util.cpp | 9 ++++----- .../samples/common/hsa_base_util.h | 15 ++++++++++---- 3 files changed, 25 insertions(+), 19 deletions(-) diff --git a/projects/rocr-runtime/samples/common/assemble.cpp b/projects/rocr-runtime/samples/common/assemble.cpp index f8b4ba8b8e..2e1827299c 100644 --- a/projects/rocr-runtime/samples/common/assemble.cpp +++ b/projects/rocr-runtime/samples/common/assemble.cpp @@ -16,7 +16,7 @@ #include "HSAILBrigObjectFile.h" namespace { - std::unordered_map mod2con; + std::unordered_map 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_EXT_STATUS_ERROR_INVALID_MODULE); } auto insert_status = mod2con.insert( - std::make_pair( - reinterpret_cast(brig_container->getBrigModule()), + std::make_pair( + brig_container->getBrigModule(), reinterpret_cast(brig_container) ) ); assert(insert_status.second); - module->handle = reinterpret_cast(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( - reinterpret_cast(brig_container->getBrigModule()), + std::make_pair( + brig_container->getBrigModule(), reinterpret_cast(brig_container) ) ); @@ -99,14 +99,14 @@ hsa_status_t ModuleCreateFromHsailString( return HSA_STATUS_ERROR_OUT_OF_RESOURCES; } - module->handle = reinterpret_cast(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_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_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_EXT_STATUS_ERROR_INVALID_MODULE); } diff --git a/projects/rocr-runtime/samples/common/hsa_base_util.cpp b/projects/rocr-runtime/samples/common/hsa_base_util.cpp index 6aecefc824..a454302cc9 100644 --- a/projects/rocr-runtime/samples/common/hsa_base_util.cpp +++ b/projects/rocr-runtime/samples/common/hsa_base_util.cpp @@ -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); diff --git a/projects/rocr-runtime/samples/common/hsa_base_util.h b/projects/rocr-runtime/samples/common/hsa_base_util.h index da0d052679..dfc32ab6fa 100644 --- a/projects/rocr-runtime/samples/common/hsa_base_util.h +++ b/projects/rocr-runtime/samples/common/hsa_base_util.h @@ -1,5 +1,6 @@ -#ifndef __HSA_UTIL__ -#define __HSA_UTIL__ +#ifndef __HSA_BASE__ +#define __HSA_BASE__ + #include #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; };