From 9a9418ad16071a38f5e77a1bf3c7f207e818b2a2 Mon Sep 17 00:00:00 2001 From: Evgeny Date: Mon, 22 Jan 2018 17:02:46 -0600 Subject: [PATCH] fixing resources releasing --- inc/rocprofiler.h | 33 ++++++-- src/core/metrics.h | 8 +- src/core/rocprofiler.cpp | 1 + src/xml/expr.h | 24 +++++- test/ctrl/run_kernel.h | 14 ++-- test/ctrl/test.cpp | 8 +- test/ctrl/test_hsa.cpp | 38 ++++----- test/ctrl/test_hsa.h | 6 ++ test/ctrl/test_pgen.h | 45 ----------- test/ctrl/test_pmgr.cpp | 141 --------------------------------- test/ctrl/test_pmgr.h | 70 ---------------- test/ctrl/tool.cpp | 2 + test/util/hsa_rsrc_factory.cpp | 63 +++++++++++---- test/util/hsa_rsrc_factory.h | 15 ++-- 14 files changed, 145 insertions(+), 323 deletions(-) delete mode 100644 test/ctrl/test_pgen.h delete mode 100644 test/ctrl/test_pmgr.cpp delete mode 100644 test/ctrl/test_pmgr.h diff --git a/inc/rocprofiler.h b/inc/rocprofiler.h index d268d6772a..9936396c63 100644 --- a/inc/rocprofiler.h +++ b/inc/rocprofiler.h @@ -76,10 +76,11 @@ extern "C" { // Profiling info objects have profiling feature info, type, parameters and data // Also profiling data samplaes can be iterated using a callback -// Profiling feature type +// Profiling feature kind typedef enum { - ROCPROFILER_FEATURE_KIND_METRIC = 0, - ROCPROFILER_FEATURE_KIND_TRACE = 1 + ROCPROFILER_FEATURE_KIND_COUNTER = 0, + ROCPROFILER_FEATURE_KIND_METRIC = 1, + ROCPROFILER_FEATURE_KIND_TRACE = 2 } rocprofiler_feature_kind_t; // Profiling feture parameter @@ -112,15 +113,24 @@ typedef struct { }; } rocprofiler_data_t; -// Profiling feature info +// Profiling feature type typedef struct { rocprofiler_feature_kind_t kind; // feature kind - const char* name; // feature name + union { + const char* name; // feature name + struct { + const char* block; // counter block name + uint32_t event; // counter event id + } counter; + }; const rocprofiler_parameter_t* parameters; // feature parameters array uint32_t parameter_count; // feature parameters count rocprofiler_data_t data; // profiling data } rocprofiler_feature_t; +// Profiling features set type +typedef void rocprofiler_feature_set_t; + //////////////////////////////////////////////////////////////////////////////// // Profiling context // @@ -158,12 +168,23 @@ typedef struct { // Create new profiling context hsa_status_t rocprofiler_open(hsa_agent_t agent, // GPU handle - rocprofiler_feature_t* features, // [in] profiling info array + rocprofiler_feature_t* features, // [in] profiling features array uint32_t feature_count, // profiling info count rocprofiler_t** context, // [out] context object uint32_t mode, // profiling mode mask rocprofiler_properties_t* properties); // profiling properties +// Add feature to e features set +hsa_status_t rocprofiler_add_feature(const rocprofiler_feature_t* feature, // [in] + rocprofiler_feature_set_t* features_set); // [in/out] profiling features set + +// Create new profiling context +hsa_status_t rocprofiler_features_set_open(hsa_agent_t agent, // GPU handle + rocprofiler_feature_set_t* features_set, // [in] profiling features set + rocprofiler_t** context, // [out] context object + uint32_t mode, // profiling mode mask + rocprofiler_properties_t* properties); // profiling properties + // Delete profiling info hsa_status_t rocprofiler_close(rocprofiler_t* context); // [in] profiling context diff --git a/src/core/metrics.h b/src/core/metrics.h index ae27fcaf82..3ff6e31c18 100644 --- a/src/core/metrics.h +++ b/src/core/metrics.h @@ -100,9 +100,11 @@ class MetricsDict { } static void Destroy() { - for (auto& entry : *map_) delete entry.second; - delete map_; - map_ = NULL; + if (map_ != NULL) { + for (auto& entry : *map_) delete entry.second; + delete map_; + map_ = NULL; + } } const Metric* Get(const std::string& name) const { diff --git a/src/core/rocprofiler.cpp b/src/core/rocprofiler.cpp index 4afbf0747e..7c5073465f 100644 --- a/src/core/rocprofiler.cpp +++ b/src/core/rocprofiler.cpp @@ -101,6 +101,7 @@ CONSTRUCTOR_API void constructor() { } DESTRUCTOR_API void destructor() { + rocprofiler::MetricsDict::Destroy(); util::HsaRsrcFactory::Destroy(); util::Logger::Destroy(); } diff --git a/src/xml/expr.h b/src/xml/expr.h index 385975975a..ca4784871c 100644 --- a/src/xml/expr.h +++ b/src/xml/expr.h @@ -18,6 +18,11 @@ class exception_t : public std::exception { const std::string str_; }; +class div_zero_exception_t : public exception_t { + public: + explicit div_zero_exception_t(const std::string& msg) : exception_t("Divide by zero exception " + msg) {} +}; + typedef uint64_t args_t; typedef std::map args_map_t; class Expr; @@ -89,7 +94,17 @@ class Expr { const expr_cache_t* GetCache() const { return cache_; } const bin_expr_t* GetTree() const { return tree_; } - args_t Eval(const args_cache_t& args) const { return tree_->Eval(args); } + args_t Eval(const args_cache_t& args) const { + args_t result = 0; + try { + result = tree_->Eval(args); + } catch (const div_zero_exception_t& e) { + if (div_zero_exc_on) std::cout << "Expr::Eval() exc(" << e.what() << ") : " << String() << std::endl; + } catch (const exception_t& e) { + throw e; + } + return result; + } std::string Lookup(const std::string& str) const { std::string result; @@ -190,6 +205,7 @@ class Expr { const expr_cache_t* const cache_; std::vector* sub_vec_; std::vector* var_vec_; + static const bool div_zero_exc_on = false; }; class add_expr_t : public bin_expr_t { @@ -213,7 +229,11 @@ class mul_expr_t : public bin_expr_t { class div_expr_t : public bin_expr_t { public: div_expr_t(const bin_expr_t* arg1, const bin_expr_t* arg2) : bin_expr_t(arg1, arg2) {} - args_t Eval(const args_cache_t& args) const { return (arg1_->Eval(args) / arg2_->Eval(args)); } + args_t Eval(const args_cache_t& args) const { + const args_t denominator = arg2_->Eval(args); + if (denominator == 0) throw div_zero_exception_t("div_expr_t::Eval()"); + return (arg1_->Eval(args) / denominator); + } std::string Symbol() const { return "/"; } }; class const_expr_t : public bin_expr_t { diff --git a/test/ctrl/run_kernel.h b/test/ctrl/run_kernel.h index 4eed641056..a731dbe17a 100644 --- a/test/ctrl/run_kernel.h +++ b/test/ctrl/run_kernel.h @@ -31,7 +31,7 @@ OF THE POSSIBILITY OF SUCH DAMAGE. #include "ctrl/test_hsa.h" #include "util/test_assert.h" -template bool RunKernel(int argc, char* argv[]) { +template bool RunKernel(int argc, char* argv[], int count = 1) { bool ret_val = false; // Create test kernel object @@ -58,11 +58,13 @@ template bool RunKernel(int argc, char* argv[]) { } // Run test kernel - ret_val = test_aql->Run(); - if (ret_val == false) { - std::cerr << "Error in running the test kernel" << std::endl; - TEST_ASSERT(ret_val); - return false; + for (int i = 0; i < count; ++i) { + ret_val = test_aql->Run(); + if (ret_val == false) { + std::cerr << "Error in running the test kernel" << std::endl; + TEST_ASSERT(ret_val); + return false; + } } // Verify the results of the execution diff --git a/test/ctrl/test.cpp b/test/ctrl/test.cpp index c9e8e71f23..780e0d010a 100644 --- a/test/ctrl/test.cpp +++ b/test/ctrl/test.cpp @@ -7,8 +7,12 @@ #include "simple_convolution/simple_convolution.h" int main(int argc, char** argv) { + const char* kiter_s = getenv("ROCP_KITER"); + const char* diter_s = getenv("ROCP_DITER"); + const int kiter = (kiter_s != NULL) ? atol(kiter_s) : 1; + const int diter = (diter_s != NULL) ? atol(diter_s) : 1; TestHsa::HsaInstantiate(); - for (int i = 0; i < 3; ++i) RunKernel(argc, argv); - TestHsa::HsaShutdown(); + for (int i = 0; i < kiter; ++i) RunKernel(argc, argv, diter); + //TestHsa::HsaShutdown(); return 0; } diff --git a/test/ctrl/test_hsa.cpp b/test/ctrl/test_hsa.cpp index 5398d97fae..5a423d6622 100644 --- a/test/ctrl/test_hsa.cpp +++ b/test/ctrl/test_hsa.cpp @@ -69,6 +69,7 @@ HsaRsrcFactory* TestHsa::HsaInstantiate(const uint32_t agent_ind) { } void TestHsa::HsaShutdown() { + if (hsa_queue_ != NULL) hsa_queue_destroy(hsa_queue_); if (hsa_rsrc_) hsa_rsrc_->Destroy(); } @@ -121,11 +122,11 @@ bool TestHsa::Setup() { // Load and Finalize Kernel Code Descriptor char* brig_path = (char*)brig_path_obj_.c_str(); - const bool ret_val = - hsa_rsrc_->LoadAndFinalize(agent_info_, brig_path, strdup(name_.c_str()), &kernel_code_desc_); - if (ret_val == false) { + code_buf_ = + hsa_rsrc_->LoadAndFinalize(agent_info_, brig_path, name_.c_str(), &hsa_exec_, &kernel_code_desc_); + if (code_buf_ == NULL) { std::cerr << "Error in loading and finalizing Kernel" << std::endl; - return ret_val; + return false; } // Stop the timer object @@ -187,32 +188,16 @@ bool TestHsa::Run() { aql.group_segment_size = group_segment_size; aql.private_segment_size = private_segment_size; // Initialize Aql packet with handle of signal + hsa_signal_store_relaxed(hsa_signal_, 1); aql.completion_signal = hsa_signal_; - // Compute the write index of queue and copy Aql packet into it - const uint64_t que_idx = hsa_queue_load_write_index_relaxed(hsa_queue_); - const uint32_t mask = hsa_queue_->size - 1; - std::clog << "> Executing kernel: \"" << name_ << "\"" << std::endl; // Start the timer object hsa_timer_.StartTimer(dispatch_timer_idx_); - // Disable packet so that submission to HW is complete - const auto header = aql.header; - aql.header = HSA_PACKET_TYPE_INVALID << HSA_PACKET_HEADER_TYPE; - - // Copy Aql packet into queue buffer - ((hsa_kernel_dispatch_packet_t*)(hsa_queue_->base_address))[que_idx & mask] = aql; - - // After AQL packet is fully copied into queue buffer - // update packet header from invalid state to valid state - std::atomic_thread_fence(std::memory_order_release); - ((hsa_kernel_dispatch_packet_t*)(hsa_queue_->base_address))[que_idx & mask].header = header; - - // Increment the write index and ring the doorbell to dispatch the kernel. - hsa_queue_store_write_index_relaxed(hsa_queue_, (que_idx + 1)); - hsa_signal_store_relaxed(hsa_queue_->doorbell_signal, que_idx); + // Submit AQL packet to the queue + const uint64_t que_idx = hsa_rsrc_->Submit(hsa_queue_, &aql); std::clog << "> Waiting on kernel dispatch signal, que_idx=" << que_idx << std::endl; @@ -251,4 +236,9 @@ void TestHsa::PrintTime() { << std::endl; } -bool TestHsa::Cleanup() { return true; } +bool TestHsa::Cleanup() { + hsa_executable_destroy(hsa_exec_); + hsa_memory_free(code_buf_); + hsa_signal_destroy(hsa_signal_); + return true; +} diff --git a/test/ctrl/test_hsa.h b/test/ctrl/test_hsa.h index ceb33b4ca9..2b3b29d72a 100644 --- a/test/ctrl/test_hsa.h +++ b/test/ctrl/test_hsa.h @@ -47,6 +47,8 @@ class TestHsa : public TestAql { total_time_taken_ = 0; setup_time_taken_ = 0; dispatch_time_taken_ = 0; + code_buf_ = NULL; + hsa_exec_ = {}; } // Get methods for Agent Info, HAS queue, HSA Resourcse Manager @@ -120,6 +122,10 @@ class TestHsa : public TestAql { // Test kernel name std::string name_; + + // Kernel code buffer + void* code_buf_; + hsa_executable_t hsa_exec_; }; #endif // TEST_CTRL_TEST_HSA_H_ diff --git a/test/ctrl/test_pgen.h b/test/ctrl/test_pgen.h deleted file mode 100644 index 36e36bd2ef..0000000000 --- a/test/ctrl/test_pgen.h +++ /dev/null @@ -1,45 +0,0 @@ -/****************************************************************************** - -Copyright ©2013 Advanced Micro Devices, Inc. All rights reserved. - -Redistribution and use in source and binary forms, with or without modification, -are permitted provided that the following conditions are met: - -Redistributions of source code must retain the above copyright notice, this list -of conditions and the following disclaimer. - -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. - -THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND -ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED -WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. -IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, -INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, -BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, -DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF -LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE -OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED -OF THE POSSIBILITY OF SUCH DAMAGE. - -*******************************************************************************/ - -#ifndef TEST_CTRL_TEST_PGEN_H_ -#define TEST_CTRL_TEST_PGEN_H_ - -#include "ctrl/test_pmgr.h" - -// SimpleConvolution: Class implements OpenCL SimpleConvolution sample -class TestPGen : public TestPMgr { - protected: - typedef hsa_ext_amd_aql_pm4_packet_t packet_t; - - packet_t* PrePacket() { return reinterpret_cast(&pre_packet_); } - packet_t* PostPacket() { return reinterpret_cast(&post_packet_); } - - public: - explicit TestPGen(TestAql* t) : TestPMgr(t) {} -}; - -#endif // TEST_CTRL_TEST_PGEN_H_ diff --git a/test/ctrl/test_pmgr.cpp b/test/ctrl/test_pmgr.cpp deleted file mode 100644 index b8e8197985..0000000000 --- a/test/ctrl/test_pmgr.cpp +++ /dev/null @@ -1,141 +0,0 @@ -/****************************************************************************** - -Copyright ©2013 Advanced Micro Devices, Inc. All rights reserved. - -Redistribution and use in source and binary forms, with or without modification, -are permitted provided that the following conditions are met: - -Redistributions of source code must retain the above copyright notice, this list -of conditions and the following disclaimer. - -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. - -THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND -ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED -WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. -IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, -INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, -BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, -DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF -LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE -OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED -OF THE POSSIBILITY OF SUCH DAMAGE. - -*******************************************************************************/ - -#include "ctrl/test_pmgr.h" - -#include - -#include "ctrl/test_assert.h" - -bool TestPMgr::AddPacketGfx9(const packet_t* packet) { - packet_t aql_packet = *packet; - - // Compute the write index of queue and copy Aql packet into it - uint64_t que_idx = hsa_queue_load_write_index_relaxed(GetQueue()); - const uint32_t mask = GetQueue()->size - 1; - packet_t* slot = (reinterpret_cast(GetQueue()->base_address)) + (que_idx & mask); - - // Disable packet so that submission to HW is complete - const auto header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE; - aql_packet.header &= (~((1ul << HSA_PACKET_HEADER_WIDTH_TYPE) - 1)) << HSA_PACKET_HEADER_TYPE; - aql_packet.header |= HSA_PACKET_TYPE_INVALID << HSA_PACKET_HEADER_TYPE; - - // Copy Aql packet into queue buffer - *slot = aql_packet; - // After AQL packet is fully copied into queue buffer - // update packet header from invalid state to valid state - auto header_atomic_ptr = reinterpret_cast*>(&slot->header); - header_atomic_ptr->store(header, std::memory_order_release); - - // Increment the write index and ring the doorbell to dispatch the kernel. - hsa_queue_store_write_index_relaxed(GetQueue(), (que_idx + 1)); - hsa_signal_store_relaxed(GetQueue()->doorbell_signal, que_idx); - - return true; -} - -bool TestPMgr::AddPacketGfx8(const packet_t* packet) { - // Create legacy devices PM4 data - const hsa_ext_amd_aql_pm4_packet_t* aql_packet = (const hsa_ext_amd_aql_pm4_packet_t*)packet; - slot_pm4_t data; - api_->hsa_ven_amd_aqlprofile_legacy_get_pm4(aql_packet, reinterpret_cast(data.words)); - - // Compute the write index of queue and copy Aql packet into it - uint64_t que_idx = hsa_queue_load_write_index_relaxed(GetQueue()); - const uint32_t mask = GetQueue()->size - 1; - - // Copy Aql/Pm4 blob into queue buffer - packet_t* ptr = (reinterpret_cast(GetQueue()->base_address)) + (que_idx & mask); - slot_pm4_t* slot = reinterpret_cast(ptr); - for (unsigned i = 1; i < SLOT_PM4_SIZE_DW; ++i) { - slot->words[i] = data.words[i]; - } - // To maintain global order to ensure the prior copy of the packet contents is made visible - // before the header is updated. - // With in-order CP it will wait until the first packet in the blob will be valid - std::atomic* header_atomic_ptr = - reinterpret_cast*>(&slot->words[0]); - header_atomic_ptr->store(data.words[0], std::memory_order_release); - - // Increment the write index and ring the doorbell to dispatch the kernel. - que_idx += SLOT_PM4_SIZE_AQLP - 1; - hsa_queue_store_write_index_relaxed(GetQueue(), (que_idx + 1)); - hsa_signal_store_relaxed(GetQueue()->doorbell_signal, que_idx); - - return true; -} - -bool TestPMgr::AddPacket(const packet_t* packet) { - const char* agent_name = GetAgentInfo()->name; - return (strncmp(agent_name, "gfx8", 4) == 0) ? AddPacketGfx8(packet) : AddPacketGfx9(packet); -} - -bool TestPMgr::Run() { - // Build Aql Pkts - const bool active = BuildPackets(); - if (active) { - // Submit Pre-Dispatch Aql packet - AddPacket(&pre_packet_); - } - - Test()->Run(); - - if (active) { - // Set post packet completion signal - post_packet_.completion_signal = post_signal_; - - // Submit Post-Dispatch Aql packet - AddPacket(&post_packet_); - - // Wait for Post-Dispatch packet to complete - hsa_signal_wait_acquire(post_signal_, HSA_SIGNAL_CONDITION_LT, 1, (uint64_t)-1, - HSA_WAIT_STATE_BLOCKED); - - // Dumping profiling data - DumpData(); - } - - return true; -} - -bool TestPMgr::Initialize(int argc, char** argv) { - TestAql::Initialize(argc, argv); - - hsa_status_t status = HSA_STATUS_ERROR; - status = hsa_signal_create(1, 0, NULL, &post_signal_); - TEST_ASSERT(status == HSA_STATUS_SUCCESS); - api_ = HsaRsrcFactory::Instance().AqlProfileApi();; - - return true; -} - -TestPMgr::TestPMgr(TestAql* t) : TestAql(t), api_(NULL) { - memset(&pre_packet_, 0, sizeof(pre_packet_)); - memset(&post_packet_, 0, sizeof(post_packet_)); - dummy_signal_.handle = 0; - post_signal_ = dummy_signal_; -} diff --git a/test/ctrl/test_pmgr.h b/test/ctrl/test_pmgr.h deleted file mode 100644 index 7e38e39655..0000000000 --- a/test/ctrl/test_pmgr.h +++ /dev/null @@ -1,70 +0,0 @@ -/****************************************************************************** - -Copyright ©2013 Advanced Micro Devices, Inc. All rights reserved. - -Redistribution and use in source and binary forms, with or without modification, -are permitted provided that the following conditions are met: - -Redistributions of source code must retain the above copyright notice, this list -of conditions and the following disclaimer. - -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. - -THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND -ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED -WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. -IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, -INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, -BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, -DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF -LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE -OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED -OF THE POSSIBILITY OF SUCH DAMAGE. - -*******************************************************************************/ - -#ifndef TEST_CTRL_TEST_PMGR_H_ -#define TEST_CTRL_TEST_PMGR_H_ - -#include -#include -#include - -#include "ctrl/test_aql.h" - -// Class implements profiling manager -class TestPMgr : public TestAql { - public: - typedef hsa_ext_amd_aql_pm4_packet_t packet_t; - explicit TestPMgr(TestAql* t); - bool Run(); - - protected: - packet_t pre_packet_; - packet_t post_packet_; - hsa_signal_t dummy_signal_; - hsa_signal_t post_signal_; - - HsaRsrcFactory::aqlprofile_pfn_t* api_; - - virtual bool BuildPackets() { return false; } - virtual bool DumpData() { return false; } - virtual bool Initialize(int argc, char** argv); - - private: - enum { - SLOT_PM4_SIZE_DW = HSA_VEN_AMD_AQLPROFILE_LEGACY_PM4_PACKET_SIZE / sizeof(uint32_t), - SLOT_PM4_SIZE_AQLP = HSA_VEN_AMD_AQLPROFILE_LEGACY_PM4_PACKET_SIZE / sizeof(packet_t) - }; - struct slot_pm4_t { - uint32_t words[SLOT_PM4_SIZE_DW]; - }; - - bool AddPacket(const packet_t* packet); - bool AddPacketGfx8(const packet_t* packet); - bool AddPacketGfx9(const packet_t* packet); -}; - -#endif // TEST_CTRL_TEST_PMGR_H_ diff --git a/test/ctrl/tool.cpp b/test/ctrl/tool.cpp index 2137265f43..07871a529d 100644 --- a/test/ctrl/tool.cpp +++ b/test/ctrl/tool.cpp @@ -184,6 +184,8 @@ void output_results(FILE* file, const rocprofiler_feature_t* features, const uns ptr = chunk_data + off; size += chunk_size; } + free(p->data.result_bytes.ptr); + fprintf(file, "size(%lu)\n", size); if (size > p->data.result_bytes.size) { fprintf(stderr, "SQTT data size is out of the result buffer size\n"); diff --git a/test/util/hsa_rsrc_factory.cpp b/test/util/hsa_rsrc_factory.cpp index 24c1a4ac3d..6e5573b1b2 100644 --- a/test/util/hsa_rsrc_factory.cpp +++ b/test/util/hsa_rsrc_factory.cpp @@ -32,6 +32,7 @@ POSSIBILITY OF SUCH DAMAGE. #include #include +#include #include #include #include @@ -111,6 +112,7 @@ static hsa_status_t GetHsaAgentsCallback(hsa_agent_t agent, void* data) { // Constructor of the class HsaRsrcFactory::HsaRsrcFactory() { // Initialize the Hsa Runtime + printf("HSA init\n"); hsa_status_t status = hsa_init(); CHECK_STATUS("Error in hsa_init", status); @@ -130,6 +132,10 @@ HsaRsrcFactory::HsaRsrcFactory() { // Destructor of the class HsaRsrcFactory::~HsaRsrcFactory() { + for (auto p : cpu_list_) free(p); + for (auto p : gpu_list_) free(p); + + printf("HSA shutdown\n"); hsa_status_t status = hsa_shut_down(); CHECK_STATUS("Error in hsa_shut_down", status); } @@ -318,8 +324,8 @@ bool HsaRsrcFactory::TransferData(void* dest_buff, void* src_buff, uint32_t leng // // @return bool true if successful, false otherwise // -bool HsaRsrcFactory::LoadAndFinalize(AgentInfo* agent_info, const char* brig_path, - char* kernel_name, hsa_executable_symbol_t* code_desc) { +void* HsaRsrcFactory::LoadAndFinalize(AgentInfo* agent_info, const char* brig_path, + const char* kernel_name, hsa_executable_t* hsa_exec, hsa_executable_symbol_t* code_desc) { // Finalize the Hsail object into code object hsa_status_t status; hsa_code_object_t code_object; @@ -333,52 +339,52 @@ bool HsaRsrcFactory::LoadAndFinalize(AgentInfo* agent_info, const char* brig_pat if (!codeStream) { std::cerr << "Error: failed to load " << filename << std::endl; assert(false); - return false; + return NULL; } // Allocate memory to read in code object from file size_t size = std::string::size_type(codeStream.tellg()); - char* codeBuff = (char*)AllocateSysMemory(agent_info, size); - if (!codeBuff) { + char* code_buf = (char*)AllocateSysMemory(agent_info, size); + if (!code_buf) { std::cerr << "Error: failed to allocate memory for code object." << std::endl; assert(false); - return false; + return NULL; } // Read the code object into allocated memory codeStream.seekg(0, std::ios::beg); - std::copy(std::istreambuf_iterator(codeStream), std::istreambuf_iterator(), codeBuff); + std::copy(std::istreambuf_iterator(codeStream), std::istreambuf_iterator(), code_buf); // De-Serialize the code object that has been read into memory - status = hsa_code_object_deserialize(codeBuff, size, NULL, &code_object); + status = hsa_code_object_deserialize(code_buf, size, NULL, &code_object); if (status != HSA_STATUS_SUCCESS) { std::cerr << "Failed to deserialize code object" << std::endl; - return false; + if (code_buf) hsa_memory_free(code_buf); + return NULL; } // Create executable. - hsa_executable_t hsaExecutable; status = - hsa_executable_create(HSA_PROFILE_FULL, HSA_EXECUTABLE_STATE_UNFROZEN, "", &hsaExecutable); + hsa_executable_create(HSA_PROFILE_FULL, HSA_EXECUTABLE_STATE_UNFROZEN, "", hsa_exec); CHECK_STATUS("Error in creating executable object", status); // Load code object. - status = hsa_executable_load_code_object(hsaExecutable, agent_info->dev_id, code_object, ""); + status = hsa_executable_load_code_object(*hsa_exec, agent_info->dev_id, code_object, ""); CHECK_STATUS("Error in loading executable object", status); // Freeze executable. - status = hsa_executable_freeze(hsaExecutable, ""); + status = hsa_executable_freeze(*hsa_exec, ""); CHECK_STATUS("Error in freezing executable object", status); // Get symbol handle. hsa_executable_symbol_t kernelSymbol; - status = hsa_executable_get_symbol(hsaExecutable, NULL, kernel_name, agent_info->dev_id, 0, + status = hsa_executable_get_symbol(*hsa_exec, NULL, kernel_name, agent_info->dev_id, 0, &kernelSymbol); CHECK_STATUS("Error in looking up kernel symbol", status); // Update output parameter *code_desc = kernelSymbol; - return true; + return code_buf; } // Add an instance of AgentInfo representing a Hsa Gpu agent @@ -411,5 +417,32 @@ bool HsaRsrcFactory::PrintGpuAgents(const std::string& header) { return true; } +uint64_t HsaRsrcFactory::Submit(hsa_queue_t* queue, void* packet) { + const uint32_t slot_size_b = 0x40; + + // adevance command queue + const uint64_t write_idx = hsa_queue_load_write_index_relaxed(queue); + hsa_queue_store_write_index_relaxed(queue, write_idx + 1); + while ((write_idx - hsa_queue_load_read_index_relaxed(queue)) >= queue->size) { + sched_yield(); + } + + uint32_t slot_idx = (uint32_t)(write_idx % queue->size); + uint32_t* queue_slot = (uint32_t*)((uintptr_t)(queue->base_address) + (slot_idx * slot_size_b)); + uint32_t* slot_data = (uint32_t*)packet; + + // Copy buffered commands into the queue slot. + // Overwrite the AQL invalid header (first dword) last. + // This prevents the slot from being read until it's fully written. + memcpy(&queue_slot[1], &slot_data[1], slot_size_b - sizeof(uint32_t)); + std::atomic* header_atomic_ptr = reinterpret_cast*>(&queue_slot[0]); + header_atomic_ptr->store(slot_data[0], std::memory_order_release); + + // ringdoor bell + hsa_signal_store_relaxed(queue->doorbell_signal, write_idx); + + return write_idx; +} + HsaRsrcFactory* HsaRsrcFactory::instance_ = NULL; HsaRsrcFactory::mutex_t HsaRsrcFactory::mutex_; diff --git a/test/util/hsa_rsrc_factory.h b/test/util/hsa_rsrc_factory.h index 2cab71a74d..c32b83f3b5 100644 --- a/test/util/hsa_rsrc_factory.h +++ b/test/util/hsa_rsrc_factory.h @@ -196,10 +196,10 @@ class HsaRsrcFactory { // @param code_desc Handle of finalized Code Descriptor that could // be used to submit for execution // - // @return bool true if successful, false otherwise + // @return code buffer, non NULL if successful, NULL otherwise // - bool LoadAndFinalize(AgentInfo* agent_info, const char* brig_path, char* kernel_name, - hsa_executable_symbol_t* code_desc); + void* LoadAndFinalize(AgentInfo* agent_info, const char* brig_path, const char* kernel_name, + hsa_executable_t* hsa_exec, hsa_executable_symbol_t* code_desc); // Add an instance of AgentInfo representing a Hsa Gpu agent void AddAgentInfo(AgentInfo* agent_info, bool gpu); @@ -207,6 +207,9 @@ class HsaRsrcFactory { // Print the various fields of Hsa Gpu Agents bool PrintGpuAgents(const std::string& header); + // Submit AQL packet to given queue + static uint64_t Submit(hsa_queue_t* queue, void* packet); + // Return AqlProfile API table typedef hsa_ven_amd_aqlprofile_1_00_pfn_t aqlprofile_pfn_t; const aqlprofile_pfn_t* AqlProfileApi() const { return &aqlprofile_api_; } @@ -225,12 +228,6 @@ class HsaRsrcFactory { static HsaRsrcFactory* instance_; static mutex_t mutex_; - // Used to maintain a list of Hsa Queue handles - std::vector queue_list_; - - // Used to maintain a list of Hsa Signal handles - std::vector signal_list_; - // Used to maintain a list of Hsa Gpu Agent Info std::vector gpu_list_;