diff --git a/.github/workflows/formatting.yml b/.github/workflows/formatting.yml index e1891a20..caf06211 100644 --- a/.github/workflows/formatting.yml +++ b/.github/workflows/formatting.yml @@ -143,3 +143,21 @@ jobs: command: review pull_number: ${{ github.event.pull_request.number }} git_dir: '.' + + missing-new-line: + runs-on: ubuntu-22.04 + + steps: + - uses: actions/checkout@v4 + + - name: Find missing new line + shell: bash + run: | + OUTFILE=missing_newline.txt + for i in $(find source/lib source/include tests samples cmake -type f | egrep -v '\.bin$'); do VAL=$(tail -c 1 ${i}); if [ -n "${VAL}" ]; then echo "- ${i}" >> ${OUTFILE}; fi; done + if [[ -f ${OUTFILE} && $(cat ${OUTFILE} | wc -l) -gt 0 ]]; then + echo -e "\nError! Source code missing new line at end of file...\n" + echo -e "\nFiles:\n" + cat ${OUTFILE} + exit 1 + fi diff --git a/samples/pc_sampling/pcs.cpp b/samples/pc_sampling/pcs.cpp index 307f292d..24167715 100644 --- a/samples/pc_sampling/pcs.cpp +++ b/samples/pc_sampling/pcs.cpp @@ -160,11 +160,12 @@ query_avail_configs_for_agent(tool_agent_info* agent_info) { // The query operation failed, so consider the PC sampling is unsupported at the agent. // This can happen if the PC sampling service is invoked within the ROCgdb. - ss << "Querying PC sampling capabilities failed with status: " << status << std::endl; + ss << "Querying PC sampling capabilities failed with status=" << status + << " :: " << rocprofiler_get_status_string(status) << std::endl; *utils::get_output_stream() << ss.str() << std::endl; return false; } - else if(agent_info->avail_configs->size() == 0) + else if(agent_info->avail_configs->empty()) { // No available configuration at the moment, so mark the PC sampling as unsupported. return false; diff --git a/source/lib/common/utility.cpp b/source/lib/common/utility.cpp index 3b37fce5..09c130cf 100644 --- a/source/lib/common/utility.cpp +++ b/source/lib/common/utility.cpp @@ -148,4 +148,4 @@ rocprofiler_debugger_continue() { debugger_block().exchange(false); } -} \ No newline at end of file +} diff --git a/source/lib/common/utility.hpp b/source/lib/common/utility.hpp index 2553cb15..b2e0eabf 100644 --- a/source/lib/common/utility.hpp +++ b/source/lib/common/utility.hpp @@ -271,4 +271,4 @@ void rocprofiler_debugger_block(); void rocprofiler_debugger_continue(); -} \ No newline at end of file +} diff --git a/source/lib/rocprofiler-sdk-codeobj/tests/CMakeLists.txt b/source/lib/rocprofiler-sdk-codeobj/tests/CMakeLists.txt index 3181522c..a9a36a89 100644 --- a/source/lib/rocprofiler-sdk-codeobj/tests/CMakeLists.txt +++ b/source/lib/rocprofiler-sdk-codeobj/tests/CMakeLists.txt @@ -29,5 +29,5 @@ set_tests_properties(${codeobj-library-test_TESTS} PROPERTIES TIMEOUT 10 LABELS target_compile_definitions(codeobj-library-test PRIVATE -DCODEOBJ_BINARY_DIR=\"${CMAKE_CURRENT_BINARY_DIR}/\") -configure_file(smallkernel.b smallkernel.b COPYONLY) +configure_file(smallkernel.bin smallkernel.bin COPYONLY) configure_file(hipcc_output.s hipcc_output.s COPYONLY) diff --git a/source/lib/rocprofiler-sdk-codeobj/tests/codeobj_library_test.cpp b/source/lib/rocprofiler-sdk-codeobj/tests/codeobj_library_test.cpp index dc9d633e..fde0f6c8 100644 --- a/source/lib/rocprofiler-sdk-codeobj/tests/codeobj_library_test.cpp +++ b/source/lib/rocprofiler-sdk-codeobj/tests/codeobj_library_test.cpp @@ -69,7 +69,7 @@ static const std::vector& GetCodeobjContents() { static std::vector buffer = []() { - std::string filename = CODEOBJ_BINARY_DIR "smallkernel.b"; + std::string filename = CODEOBJ_BINARY_DIR "smallkernel.bin"; std::ifstream file(filename.data(), std::ios::binary); using iterator_t = std::istreambuf_iterator; @@ -141,7 +141,7 @@ TEST(codeobj_library, decoder_component) CodeobjDecoderComponent component(objdata.data(), objdata.size()); - std::string kernel_with_protocol = "file://" CODEOBJ_BINARY_DIR "smallkernel.b"; + std::string kernel_with_protocol = "file://" CODEOBJ_BINARY_DIR "smallkernel.bin"; LoadedCodeobjDecoder loadecomp(kernel_with_protocol.data(), loaded_offset, objdata.size()); ASSERT_EQ(component.m_symbol_map.size(), 1); diff --git a/source/lib/rocprofiler-sdk-codeobj/tests/smallkernel.b b/source/lib/rocprofiler-sdk-codeobj/tests/smallkernel.bin similarity index 100% rename from source/lib/rocprofiler-sdk-codeobj/tests/smallkernel.b rename to source/lib/rocprofiler-sdk-codeobj/tests/smallkernel.bin diff --git a/source/lib/rocprofiler-sdk/agent_profile.cpp b/source/lib/rocprofiler-sdk/agent_profile.cpp index 35470079..5b0105cf 100644 --- a/source/lib/rocprofiler-sdk/agent_profile.cpp +++ b/source/lib/rocprofiler-sdk/agent_profile.cpp @@ -47,4 +47,4 @@ rocprofiler_sample_agent_profile_counting_service(rocprofiler_context_id_t con return rocprofiler::counters::read_agent_ctx( rocprofiler::context::get_registered_context(context_id), user_data, flags); } -} \ No newline at end of file +} diff --git a/source/lib/rocprofiler-sdk/aql/helpers.cpp b/source/lib/rocprofiler-sdk/aql/helpers.cpp index 0111732f..c3d6101f 100644 --- a/source/lib/rocprofiler-sdk/aql/helpers.cpp +++ b/source/lib/rocprofiler-sdk/aql/helpers.cpp @@ -21,15 +21,15 @@ // SOFTWARE. #include "lib/rocprofiler-sdk/aql/helpers.hpp" - -#include - -#include - #include "lib/common/logging.hpp" #include "lib/common/synchronized.hpp" #include "lib/common/utility.hpp" #include "lib/rocprofiler-sdk/counters/id_decode.hpp" +#include "lib/rocprofiler-sdk/hsa/hsa.hpp" + +#include + +#include namespace rocprofiler { @@ -114,11 +114,15 @@ get_dim_info(rocprofiler_agent_id_t agent, } rocprofiler_status_t -set_profiler_active_on_queue(const AmdExtTable& api, - hsa_amd_memory_pool_t pool, +set_profiler_active_on_queue(hsa_amd_memory_pool_t pool, hsa_agent_t hsa_agent, const rocprofiler_profile_pkt_cb& packet_submit) { + CHECK(hsa::get_amd_ext_table() != nullptr); + CHECK(hsa::get_amd_ext_table()->hsa_amd_memory_pool_allocate_fn != nullptr); + CHECK(hsa::get_amd_ext_table()->hsa_amd_agents_allow_access_fn != nullptr); + CHECK(hsa::get_amd_ext_table()->hsa_amd_memory_pool_free_fn != nullptr); + // Inject packet to enable profiling of other process queues on this queue hsa_ven_amd_aqlprofile_profile_t profile{}; profile.agent = hsa_agent; @@ -134,15 +138,15 @@ set_profiler_active_on_queue(const AmdExtTable& api, const size_t mask = 0x1000 - 1; auto size = (profile.command_buffer.size + mask) & ~mask; - if(api.hsa_amd_memory_pool_allocate_fn(pool, size, 0, &profile.command_buffer.ptr) != - HSA_STATUS_SUCCESS) + if(hsa::get_amd_ext_table()->hsa_amd_memory_pool_allocate_fn( + pool, size, 0, &profile.command_buffer.ptr) != HSA_STATUS_SUCCESS) { ROCP_WARNING << "Failed to allocate memory to enable profile command on agent, some " "counters will be unavailable"; return ROCPROFILER_STATUS_ERROR; } - if(api.hsa_amd_agents_allow_access_fn(1, &hsa_agent, nullptr, profile.command_buffer.ptr) != - HSA_STATUS_SUCCESS) + if(hsa::get_amd_ext_table()->hsa_amd_agents_allow_access_fn( + 1, &hsa_agent, nullptr, profile.command_buffer.ptr) != HSA_STATUS_SUCCESS) { ROCP_WARNING << "Agent cannot access memory, some counters will be unavailable"; return ROCPROFILER_STATUS_ERROR; @@ -157,7 +161,7 @@ set_profiler_active_on_queue(const AmdExtTable& api, } packet_submit(packet); - api.hsa_amd_memory_pool_free_fn(profile.command_buffer.ptr); + hsa::get_amd_ext_table()->hsa_amd_memory_pool_free_fn(profile.command_buffer.ptr); return ROCPROFILER_STATUS_SUCCESS; } diff --git a/source/lib/rocprofiler-sdk/aql/helpers.hpp b/source/lib/rocprofiler-sdk/aql/helpers.hpp index 802cede6..371f35a0 100644 --- a/source/lib/rocprofiler-sdk/aql/helpers.hpp +++ b/source/lib/rocprofiler-sdk/aql/helpers.hpp @@ -22,18 +22,18 @@ #pragma once -#include -#include -#include - -#include - -#include - #include "lib/rocprofiler-sdk/agent.hpp" #include "lib/rocprofiler-sdk/counters/metrics.hpp" #include "lib/rocprofiler-sdk/hsa/rocprofiler_packet.hpp" +#include + +#include + +#include +#include +#include + namespace rocprofiler { namespace aql @@ -62,8 +62,7 @@ set_dim_id_from_sample(rocprofiler_counter_instance_id_t& id, size_t sample_id); rocprofiler_status_t -set_profiler_active_on_queue(const AmdExtTable& api, - hsa_amd_memory_pool_t pool, +set_profiler_active_on_queue(hsa_amd_memory_pool_t pool, hsa_agent_t hsa_agent, const rocprofiler_profile_pkt_cb& packet_submit); } // namespace aql diff --git a/source/lib/rocprofiler-sdk/aql/tests/aql_test.cpp b/source/lib/rocprofiler-sdk/aql/tests/aql_test.cpp index 242bd02f..f633da2c 100644 --- a/source/lib/rocprofiler-sdk/aql/tests/aql_test.cpp +++ b/source/lib/rocprofiler-sdk/aql/tests/aql_test.cpp @@ -219,4 +219,4 @@ TEST(aql_profile, test_aql_packet) // Why is this valid? TestAqlPacket test_pkt2(false); } -*/ \ No newline at end of file +*/ diff --git a/source/lib/rocprofiler-sdk/counters/agent_profiling.cpp b/source/lib/rocprofiler-sdk/counters/agent_profiling.cpp index 13368abc..2a4c9a0c 100644 --- a/source/lib/rocprofiler-sdk/counters/agent_profiling.cpp +++ b/source/lib/rocprofiler-sdk/counters/agent_profiling.cpp @@ -21,18 +21,21 @@ // SOFTWARE. #include "lib/rocprofiler-sdk/counters/agent_profiling.hpp" -#include -#include - #include "lib/common/logging.hpp" #include "lib/rocprofiler-sdk/buffer.hpp" #include "lib/rocprofiler-sdk/context/context.hpp" #include "lib/rocprofiler-sdk/counters/controller.hpp" #include "lib/rocprofiler-sdk/counters/core.hpp" #include "lib/rocprofiler-sdk/hsa/agent_cache.hpp" +#include "lib/rocprofiler-sdk/hsa/hsa.hpp" #include "lib/rocprofiler-sdk/hsa/queue_controller.hpp" #include "lib/rocprofiler-sdk/hsa/rocprofiler_packet.hpp" -#include "rocprofiler-sdk/fwd.h" + +#include + +#include +#include +#include namespace rocprofiler { @@ -46,13 +49,15 @@ hsa_inited() } uint64_t -submitPacket(const CoreApiTable& table, hsa_queue_t* queue, const void* packet) +submitPacket(hsa_queue_t* queue, const void* packet) { const uint32_t pkt_size = 0x40; // advance command queue - const uint64_t write_idx = table.hsa_queue_add_write_index_scacq_screl_fn(queue, 1); - while((write_idx - table.hsa_queue_load_read_index_relaxed_fn(queue)) >= queue->size) + const uint64_t write_idx = + hsa::get_core_table()->hsa_queue_add_write_index_scacq_screl_fn(queue, 1); + while((write_idx - hsa::get_core_table()->hsa_queue_load_read_index_relaxed_fn(queue)) >= + queue->size) { sched_yield(); } @@ -74,7 +79,7 @@ submitPacket(const CoreApiTable& table, hsa_queue_t* queue, const void* packet) header_atomic_ptr->store(slot_data[0], std::memory_order_release); // ringdoor bell - table.hsa_signal_store_relaxed_fn(queue->doorbell_signal, write_idx); + hsa::get_core_table()->hsa_signal_store_relaxed_fn(queue->doorbell_signal, write_idx); return write_idx; } @@ -148,7 +153,7 @@ agent_async_handler(hsa_signal_value_t /*signal_v*/, void* data) } // reset the signal to allow another sample to start - callback_data.table.hsa_signal_store_relaxed_fn(callback_data.completion, 1); + hsa::get_core_table()->hsa_signal_store_relaxed_fn(callback_data.completion, 1); return true; } @@ -167,31 +172,36 @@ init_callback_data(rocprofiler::counters::agent_callback_data& callback_data, callback_data.packet = construct_aql_pkt(callback_data.profile); callback_data.queue = agent.profile_queue(); - callback_data.table = CHECK_NOTNULL(hsa::get_queue_controller())->get_core_table(); if(callback_data.completion.handle != 0) return; + CHECK(hsa::get_core_table() != nullptr); + CHECK(hsa::get_amd_ext_table() != nullptr); + CHECK(hsa::get_core_table()->hsa_signal_create_fn != nullptr); + CHECK(hsa::get_core_table()->hsa_signal_wait_relaxed_fn != nullptr); + CHECK(hsa::get_core_table()->hsa_signal_store_relaxed_fn != nullptr); + CHECK(hsa::get_amd_ext_table()->hsa_amd_signal_async_handler_fn != nullptr); + // Tri-state signal // 1: allow next sample to start // 0: sample in progress // -1: sample complete - CHECK_EQ(callback_data.table.hsa_signal_create_fn(1, 0, nullptr, &callback_data.completion), + CHECK_EQ(hsa::get_core_table()->hsa_signal_create_fn(1, 0, nullptr, &callback_data.completion), HSA_STATUS_SUCCESS); // Signal to manage the startup of the context. Allows us to ensure that // the AQL packet we inject with start_context() completes before returning - CHECK_EQ(callback_data.table.hsa_signal_create_fn(1, 0, nullptr, &callback_data.start_signal), - HSA_STATUS_SUCCESS); + CHECK_EQ( + hsa::get_core_table()->hsa_signal_create_fn(1, 0, nullptr, &callback_data.start_signal), + HSA_STATUS_SUCCESS); // Setup callback // NOLINTBEGIN(performance-no-int-to-ptr) - CHECK_EQ(CHECK_NOTNULL(hsa::get_queue_controller()) - ->get_ext_table() - .hsa_amd_signal_async_handler_fn(callback_data.completion, - HSA_SIGNAL_CONDITION_LT, - 0, - agent_async_handler, - (void*) &callback_data), + CHECK_EQ(hsa::get_amd_ext_table()->hsa_amd_signal_async_handler_fn(callback_data.completion, + HSA_SIGNAL_CONDITION_LT, + 0, + agent_async_handler, + &callback_data), HSA_STATUS_SUCCESS); // NOLINTEND(performance-no-int-to-ptr) @@ -207,23 +217,20 @@ init_callback_data(rocprofiler::counters::agent_callback_data& callback_data, CHECK(agent.get_hsa_agent().handle != 0); aql::set_profiler_active_on_queue( - CHECK_NOTNULL(hsa::get_queue_controller())->get_ext_table(), - agent.cpu_pool(), - agent.get_hsa_agent(), - [&](hsa::rocprofiler_packet pkt) { + agent.cpu_pool(), agent.get_hsa_agent(), [&](hsa::rocprofiler_packet pkt) { pkt.ext_amd_aql_pm4.completion_signal = callback_data.completion; - submitPacket(callback_data.table, callback_data.queue, (void*) &pkt); + submitPacket(callback_data.queue, (void*) &pkt); constexpr auto timeout_hint = std::chrono::duration_cast(std::chrono::seconds{1}); - if(callback_data.table.hsa_signal_wait_relaxed_fn(callback_data.completion, - HSA_SIGNAL_CONDITION_EQ, - 0, - timeout_hint.count(), - HSA_WAIT_STATE_ACTIVE) != 0) + if(hsa::get_core_table()->hsa_signal_wait_relaxed_fn(callback_data.completion, + HSA_SIGNAL_CONDITION_EQ, + 0, + timeout_hint.count(), + HSA_WAIT_STATE_ACTIVE) != 0) { ROCP_FATAL << "Could not set agent to be profiled"; } - callback_data.table.hsa_signal_store_relaxed_fn(callback_data.completion, 1); + hsa::get_core_table()->hsa_signal_store_relaxed_fn(callback_data.completion, 1); }); } } // namespace @@ -284,16 +291,16 @@ read_agent_ctx(const context::context* ctx, if(callback_data.profile->reqired_hw_counters.empty()) { callback_data.user_data = user_data; - callback_data.table.hsa_signal_store_relaxed_fn(callback_data.completion, -1); + hsa::get_core_table()->hsa_signal_store_relaxed_fn(callback_data.completion, -1); // Wait for the barrier/read packet to complete if(flags != ROCPROFILER_COUNTER_FLAG_ASYNC) { // Wait for any inprogress samples to complete before returning - callback_data.table.hsa_signal_wait_relaxed_fn(callback_data.completion, - HSA_SIGNAL_CONDITION_EQ, - 1, - UINT64_MAX, - HSA_WAIT_STATE_ACTIVE); + hsa::get_core_table()->hsa_signal_wait_relaxed_fn(callback_data.completion, + HSA_SIGNAL_CONDITION_EQ, + 1, + UINT64_MAX, + HSA_WAIT_STATE_ACTIVE); } continue; } @@ -307,28 +314,26 @@ read_agent_ctx(const context::context* ctx, agent->get_rocp_agent()->simd_arrays_per_engine); // Submit the read packet to the queue - submitPacket(callback_data.table, - agent->profile_queue(), - (void*) &callback_data.packet->packets.read_packet); + submitPacket(agent->profile_queue(), &callback_data.packet->packets.read_packet); // Submit a barrier packet. This is needed to flush hardware caches. Without this // the read packet may not have the correct data. rocprofiler::hsa::rocprofiler_packet barrier{}; barrier.barrier_and.header = header_pkt(HSA_PACKET_TYPE_BARRIER_AND); barrier.barrier_and.completion_signal = callback_data.completion; - callback_data.table.hsa_signal_store_relaxed_fn(callback_data.completion, 0); + hsa::get_core_table()->hsa_signal_store_relaxed_fn(callback_data.completion, 0); callback_data.user_data = user_data; - submitPacket(callback_data.table, agent->profile_queue(), (void*) &barrier.barrier_and); + submitPacket(agent->profile_queue(), &barrier.barrier_and); // Wait for the barrier/read packet to complete if(flags != ROCPROFILER_COUNTER_FLAG_ASYNC) { // Wait for any inprogress samples to complete before returning - callback_data.table.hsa_signal_wait_relaxed_fn(callback_data.completion, - HSA_SIGNAL_CONDITION_EQ, - 1, - UINT64_MAX, - HSA_WAIT_STATE_ACTIVE); + hsa::get_core_table()->hsa_signal_wait_relaxed_fn(callback_data.completion, + HSA_SIGNAL_CONDITION_EQ, + 1, + UINT64_MAX, + HSA_WAIT_STATE_ACTIVE); } } @@ -458,17 +463,15 @@ start_agent_ctx(const context::context* ctx) } callback_data.packet->packets.start_packet.completion_signal = callback_data.start_signal; - callback_data.table.hsa_signal_store_relaxed_fn(callback_data.start_signal, 1); - submitPacket(callback_data.table, - agent->profile_queue(), - (void*) &callback_data.packet->packets.start_packet); + hsa::get_core_table()->hsa_signal_store_relaxed_fn(callback_data.start_signal, 1); + submitPacket(agent->profile_queue(), &callback_data.packet->packets.start_packet); // Wait for startup to finish before continuing - callback_data.table.hsa_signal_wait_relaxed_fn(callback_data.start_signal, - HSA_SIGNAL_CONDITION_EQ, - 0, - UINT64_MAX, - HSA_WAIT_STATE_ACTIVE); + hsa::get_core_table()->hsa_signal_wait_relaxed_fn(callback_data.start_signal, + HSA_SIGNAL_CONDITION_EQ, + 0, + UINT64_MAX, + HSA_WAIT_STATE_ACTIVE); } agent_ctx.status.exchange( @@ -517,17 +520,15 @@ stop_agent_ctx(const context::context* ctx) if(!callback_data.profile->reqired_hw_counters.empty()) { // Remove when AQL is updated to not require stop to be called first - submitPacket(callback_data.table, - agent->profile_queue(), - (void*) &callback_data.packet->packets.stop_packet); + submitPacket(agent->profile_queue(), &callback_data.packet->packets.stop_packet); } // Wait for the stop packet to complete - callback_data.table.hsa_signal_wait_relaxed_fn(callback_data.completion, - HSA_SIGNAL_CONDITION_EQ, - 1, - UINT64_MAX, - HSA_WAIT_STATE_ACTIVE); + hsa::get_core_table()->hsa_signal_wait_relaxed_fn(callback_data.completion, + HSA_SIGNAL_CONDITION_EQ, + 1, + UINT64_MAX, + HSA_WAIT_STATE_ACTIVE); } agent_ctx.status.exchange( @@ -553,7 +554,7 @@ agent_profile_hsa_registration() agent_callback_data::~agent_callback_data() { - if(completion.handle != 0) table.hsa_signal_destroy_fn(completion); + if(completion.handle != 0) hsa::get_core_table()->hsa_signal_destroy_fn(completion); } } // namespace counters } // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk/counters/agent_profiling.hpp b/source/lib/rocprofiler-sdk/counters/agent_profiling.hpp index 4bbcf8a9..be2bec04 100644 --- a/source/lib/rocprofiler-sdk/counters/agent_profiling.hpp +++ b/source/lib/rocprofiler-sdk/counters/agent_profiling.hpp @@ -21,13 +21,13 @@ // SOFTWARE. #pragma once +#include "lib/rocprofiler-sdk/counters/controller.hpp" +#include "lib/rocprofiler-sdk/hsa/aql_packet.hpp" + #include #include #include -#include "lib/rocprofiler-sdk/counters/controller.hpp" -#include "lib/rocprofiler-sdk/hsa/aql_packet.hpp" - namespace rocprofiler { namespace context @@ -39,41 +39,40 @@ namespace counters { struct agent_callback_data { - uint64_t context_idx{0}; - CoreApiTable table; - hsa_queue_t* queue{nullptr}; - std::unique_ptr packet; + uint64_t context_idx = 0; + hsa_queue_t* queue = nullptr; + std::unique_ptr packet = {}; // Tri-state signal used to know what the current state of processing // a sample is. The states are: // 1: allow next sample to start (i.e. no in progress work) // 0: sample in progress // -1: sample complete (i.e. signal for caller that sample is ready) - hsa_signal_t completion{.handle = 0}; + hsa_signal_t completion = {.handle = 0}; + hsa_signal_t start_signal = {.handle = 0}; + rocprofiler_user_data_t user_data = {.value = 0}; + rocprofiler_user_data_t callback_data = {.value = 0}; - hsa_signal_t start_signal{.handle = 0}; - rocprofiler_user_data_t user_data{.value = 0}; - rocprofiler_user_data_t callback_data{.ptr = nullptr}; - - std::shared_ptr profile; - rocprofiler_agent_id_t agent_id{.handle = 0}; - rocprofiler_agent_profile_callback_t cb; - rocprofiler_buffer_id_t buffer{.handle = 0}; - bool set_profile{false}; + std::shared_ptr profile = {}; + rocprofiler_agent_id_t agent_id = {.handle = 0}; + rocprofiler_agent_profile_callback_t cb = nullptr; + rocprofiler_buffer_id_t buffer = {.handle = 0}; + bool set_profile = false; agent_callback_data() = default; - agent_callback_data(agent_callback_data&& other) noexcept - : table(other.table) - , queue(other.queue) - , packet(std::move(other.packet)) - , completion(other.completion) - , start_signal(other.start_signal) - , user_data(other.user_data) - , callback_data(other.callback_data) - , profile(other.profile) - , agent_id(other.agent_id) - , cb(other.cb) - , buffer(other.buffer){}; + agent_callback_data(agent_callback_data&& rhs) noexcept + : queue(rhs.queue) + , packet(std::move(rhs.packet)) + , completion(rhs.completion) + , start_signal(rhs.start_signal) + , user_data(rhs.user_data) + , callback_data(rhs.callback_data) + , profile(rhs.profile) + , agent_id(rhs.agent_id) + , cb(rhs.cb) + , buffer(rhs.buffer) + {} + ~agent_callback_data(); }; @@ -108,7 +107,7 @@ read_agent_ctx(const context::context* ctx, rocprofiler_counter_flag_t flags); uint64_t -submitPacket(const CoreApiTable& table, hsa_queue_t* queue, const void* packet); +submitPacket(hsa_queue_t* queue, const void* packet); } // namespace counters -} // namespace rocprofiler \ No newline at end of file +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk/counters/dimensions.hpp b/source/lib/rocprofiler-sdk/counters/dimensions.hpp index ee88d032..6bbb996c 100644 --- a/source/lib/rocprofiler-sdk/counters/dimensions.hpp +++ b/source/lib/rocprofiler-sdk/counters/dimensions.hpp @@ -85,4 +85,4 @@ struct formatter return fmt::format_to(ctx.out(), "[{}, {}]", dims.name(), dims.size()); } }; -} // namespace fmt \ No newline at end of file +} // namespace fmt diff --git a/source/lib/rocprofiler-sdk/counters/dispatch_handlers.cpp b/source/lib/rocprofiler-sdk/counters/dispatch_handlers.cpp index 8ee7c994..03ae4020 100644 --- a/source/lib/rocprofiler-sdk/counters/dispatch_handlers.cpp +++ b/source/lib/rocprofiler-sdk/counters/dispatch_handlers.cpp @@ -277,4 +277,4 @@ completed_cb(const context::context* ctx, } } } // namespace counters -} // namespace rocprofiler \ No newline at end of file +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk/counters/dispatch_handlers.hpp b/source/lib/rocprofiler-sdk/counters/dispatch_handlers.hpp index 894091dd..918ff98b 100644 --- a/source/lib/rocprofiler-sdk/counters/dispatch_handlers.hpp +++ b/source/lib/rocprofiler-sdk/counters/dispatch_handlers.hpp @@ -1,4 +1,3 @@ - // MIT License // // Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. @@ -54,4 +53,4 @@ queue_cb(const context::context* ctx, const context::correlation_id* correlation_id); } // namespace counters -} // namespace rocprofiler \ No newline at end of file +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk/counters/tests/agent_profiling.cpp b/source/lib/rocprofiler-sdk/counters/tests/agent_profiling.cpp index a995d8a0..8af74cfa 100644 --- a/source/lib/rocprofiler-sdk/counters/tests/agent_profiling.cpp +++ b/source/lib/rocprofiler-sdk/counters/tests/agent_profiling.cpp @@ -31,8 +31,8 @@ #include "lib/rocprofiler-sdk/hsa/agent_cache.hpp" #include "lib/rocprofiler-sdk/hsa/queue_controller.hpp" #include "lib/rocprofiler-sdk/registration.hpp" -#include "rocprofiler-sdk/buffer.h" +#include #include #include #include @@ -51,7 +51,6 @@ using namespace rocprofiler::counters::test_constants; using namespace rocprofiler::counters::testing; -using namespace rocprofiler::counters; using namespace rocprofiler; #define ROCPROFILER_CALL(result, msg) \ @@ -76,17 +75,17 @@ auto findDeviceMetrics(const hsa::AgentCache& agent, const std::unordered_set& metrics) { std::vector ret; - auto all_counters = counters::getMetricMap(); + const auto* all_counters = counters::getMetricMap(); ROCP_ERROR << "Looking up counters for " << std::string(agent.name()); - auto gfx_metrics = common::get_val(*all_counters, std::string(agent.name())); + const auto* gfx_metrics = common::get_val(*all_counters, std::string(agent.name())); if(!gfx_metrics) { ROCP_ERROR << "No counters found for " << std::string(agent.name()); return ret; } - for(auto& counter : *gfx_metrics) + for(const auto& counter : *gfx_metrics) { if(metrics.count(counter.name()) > 0 || metrics.empty()) { @@ -102,6 +101,8 @@ test_init() HsaApiTable table; table.amd_ext_ = &get_ext_table(); table.core_ = &get_api_table(); + rocprofiler::hsa::copy_table(table.core_, 0); + rocprofiler::hsa::copy_table(table.amd_ext_, 0); agent::construct_agent_cache(&table); ASSERT_TRUE(hsa::get_queue_controller() != nullptr); hsa::get_queue_controller()->init(get_api_table(), get_ext_table()); @@ -292,10 +293,7 @@ class agent_profile_test : public ::testing::Test // Set state of the queue to allow profiling (may not be needed since AQL // may do this in the future). aql::set_profiler_active_on_queue( - get_ext_table(), - agent.cpu_pool(), - agent.get_hsa_agent(), - [&](hsa::rocprofiler_packet pkt) { + agent.cpu_pool(), agent.get_hsa_agent(), [&](hsa::rocprofiler_packet pkt) { pkt.ext_amd_aql_pm4.completion_signal = completion_signal; submitPacket(queue, (void*) &pkt); @@ -477,4 +475,4 @@ TEST_F(agent_profile_test, sync_sq_waves_verify) ROCP_ERROR << fmt::format("Name: {} Counter value: {}", info.name, val.counter_value); EXPECT_GT(val.counter_value, 0.0); } -} \ No newline at end of file +} diff --git a/source/lib/rocprofiler-sdk/counters/tests/agent_profiling.hpp b/source/lib/rocprofiler-sdk/counters/tests/agent_profiling.hpp index 7c3d524c..5cf388ee 100644 --- a/source/lib/rocprofiler-sdk/counters/tests/agent_profiling.hpp +++ b/source/lib/rocprofiler-sdk/counters/tests/agent_profiling.hpp @@ -20,4 +20,4 @@ // OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE // SOFTWARE. -#pragma once \ No newline at end of file +#pragma once diff --git a/source/lib/rocprofiler-sdk/counters/tests/code_object_loader.cpp b/source/lib/rocprofiler-sdk/counters/tests/code_object_loader.cpp index 27162a82..6d9eb13a 100644 --- a/source/lib/rocprofiler-sdk/counters/tests/code_object_loader.cpp +++ b/source/lib/rocprofiler-sdk/counters/tests/code_object_loader.cpp @@ -103,4 +103,4 @@ search_hasco(const common::filesystem::path& directory, std::string& filename) } } // namespace testing } // namespace counters -} // namespace rocprofiler \ No newline at end of file +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk/counters/tests/hsa_tables.cpp b/source/lib/rocprofiler-sdk/counters/tests/hsa_tables.cpp index f29ce439..83c7cb70 100644 --- a/source/lib/rocprofiler-sdk/counters/tests/hsa_tables.cpp +++ b/source/lib/rocprofiler-sdk/counters/tests/hsa_tables.cpp @@ -19,12 +19,14 @@ // 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 IN THE // SOFTWARE. + #include #include #include #include #include +#include #include namespace rocprofiler @@ -38,6 +40,9 @@ get_ext_table() { static auto _v = []() { auto val = AmdExtTable{}; + val.version.major_id = HSA_AMD_EXT_API_TABLE_MAJOR_VERSION; + val.version.minor_id = sizeof(AmdExtTable); + val.version.step_id = HSA_AMD_EXT_API_TABLE_STEP_VERSION; val.hsa_amd_coherency_get_type_fn = hsa_amd_coherency_get_type; val.hsa_amd_coherency_set_type_fn = hsa_amd_coherency_set_type; val.hsa_amd_profiling_set_profiler_enabled_fn = hsa_amd_profiling_set_profiler_enabled; @@ -122,6 +127,9 @@ get_api_table() { static auto _v = []() { auto val = CoreApiTable{}; + val.version.major_id = HSA_CORE_API_TABLE_MAJOR_VERSION; + val.version.minor_id = sizeof(CoreApiTable); + val.version.step_id = HSA_CORE_API_TABLE_STEP_VERSION; val.hsa_init_fn = hsa_init; val.hsa_shut_down_fn = hsa_shut_down; val.hsa_system_get_info_fn = hsa_system_get_info; diff --git a/source/lib/rocprofiler-sdk/counters/tests/hsa_tables.hpp b/source/lib/rocprofiler-sdk/counters/tests/hsa_tables.hpp index 92b3a687..b6a6c400 100644 --- a/source/lib/rocprofiler-sdk/counters/tests/hsa_tables.hpp +++ b/source/lib/rocprofiler-sdk/counters/tests/hsa_tables.hpp @@ -38,4 +38,4 @@ CoreApiTable& get_api_table(); } // namespace test_constants } // namespace counters -} // namespace rocprofiler \ No newline at end of file +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk/external_correlation.cpp b/source/lib/rocprofiler-sdk/external_correlation.cpp index 8646f9f5..2502443f 100644 --- a/source/lib/rocprofiler-sdk/external_correlation.cpp +++ b/source/lib/rocprofiler-sdk/external_correlation.cpp @@ -146,10 +146,10 @@ external_correlation::pop(rocprofiler_thread_id_t tid) { static auto default_tid = get_default_tid(); - return data.wlock( - [](external_correlation_map_t& _data, rocprofiler_thread_id_t tid_v) { + return data.rlock( + [](const external_correlation_map_t& _data, rocprofiler_thread_id_t tid_v) { if(_data.count(tid_v) == 0) return empty_user_data; - auto& itr = _data.at(tid_v); + const auto& itr = _data.at(tid_v); return itr.wlock([tid_v](external_correlation_stack_t& data_stack) { if(data_stack.empty()) return empty_user_data; auto ret = data_stack.back(); diff --git a/source/lib/rocprofiler-sdk/hsa/queue.cpp b/source/lib/rocprofiler-sdk/hsa/queue.cpp index 5e8529b3..7bfaeb29 100644 --- a/source/lib/rocprofiler-sdk/hsa/queue.cpp +++ b/source/lib/rocprofiler-sdk/hsa/queue.cpp @@ -510,16 +510,11 @@ Queue::Queue(const AgentCache& agent, CHECK(_agent.get_hsa_agent().handle != 0); // Set state of the queue to allow profiling aql::set_profiler_active_on_queue( - CHECK_NOTNULL(hsa::get_queue_controller())->get_ext_table(), - _agent.cpu_pool(), - _agent.get_hsa_agent(), - [&](hsa::rocprofiler_packet pkt) { + _agent.cpu_pool(), _agent.get_hsa_agent(), [&](hsa::rocprofiler_packet pkt) { hsa_signal_t completion; create_signal(0, &completion); pkt.ext_amd_aql_pm4.completion_signal = completion; - counters::submitPacket(CHECK_NOTNULL(hsa::get_queue_controller())->get_core_table(), - _intercept_queue, - &pkt); + counters::submitPacket(_intercept_queue, &pkt); constexpr auto timeout_hint = std::chrono::duration_cast(std::chrono::seconds{1}); if(core_api.hsa_signal_wait_relaxed_fn(completion, diff --git a/source/lib/rocprofiler-sdk/pc_sampling/parser/pc_record_interface.cpp b/source/lib/rocprofiler-sdk/pc_sampling/parser/pc_record_interface.cpp index abaaac84..278866ac 100644 --- a/source/lib/rocprofiler-sdk/pc_sampling/parser/pc_record_interface.cpp +++ b/source/lib/rocprofiler-sdk/pc_sampling/parser/pc_record_interface.cpp @@ -114,4 +114,4 @@ PCSamplingParserContext::generate_upcoming_pc_record( buff->emplace(ROCPROFILER_BUFFER_CATEGORY_PC_SAMPLING, ROCPROFILER_PC_SAMPLING_RECORD_SAMPLE, samples[i]); -}; \ No newline at end of file +} diff --git a/source/lib/rocprofiler-sdk/pc_sampling/tests/pc_sampling_internals.hpp b/source/lib/rocprofiler-sdk/pc_sampling/tests/pc_sampling_internals.hpp index a752c249..d7e532c1 100644 --- a/source/lib/rocprofiler-sdk/pc_sampling/tests/pc_sampling_internals.hpp +++ b/source/lib/rocprofiler-sdk/pc_sampling/tests/pc_sampling_internals.hpp @@ -60,4 +60,4 @@ get_active_pc_sampling_service(); } // namespace hsa } // namespace pc_sampling -} // namespace rocprofiler \ No newline at end of file +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk/rocprofiler.cpp b/source/lib/rocprofiler-sdk/rocprofiler.cpp index 37e85fba..8ab8274b 100644 --- a/source/lib/rocprofiler-sdk/rocprofiler.cpp +++ b/source/lib/rocprofiler-sdk/rocprofiler.cpp @@ -108,9 +108,8 @@ ROCPROFILER_STATUS_STRING(ROCPROFILER_STATUS_ERROR_NO_HARDWARE_COUNTERS, ROCPROFILER_STATUS_STRING(ROCPROFILER_STATUS_ERROR_AGENT_MISMATCH, "Counter profile agent does not match the agent in the context") ROCPROFILER_STATUS_STRING(ROCPROFILER_STATUS_ERROR_NOT_AVAILABLE, - "The service is not available." - "Please refer to API functions that return this status code" - "for more information.") + "The service is not available. Please refer to API functions that return " + "this status code for more information.") ROCPROFILER_STATUS_STRING(ROCPROFILER_STATUS_ERROR_EXCEEDS_HW_LIMIT, "Request exceeds the capabilities of the hardware to collect") diff --git a/source/lib/rocprofiler-sdk/thread_trace/tests/att_packet_test.cpp b/source/lib/rocprofiler-sdk/thread_trace/tests/att_packet_test.cpp index 129a71cd..d65430c9 100644 --- a/source/lib/rocprofiler-sdk/thread_trace/tests/att_packet_test.cpp +++ b/source/lib/rocprofiler-sdk/thread_trace/tests/att_packet_test.cpp @@ -19,6 +19,7 @@ // 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 IN THE // SOFTWARE. + #include "lib/rocprofiler-sdk/agent.hpp" #include "lib/rocprofiler-sdk/aql/helpers.hpp" #include "lib/rocprofiler-sdk/aql/packet_construct.hpp" @@ -251,4 +252,4 @@ TEST(thread_trace, perfcounters_aql_options_test) sqtt_default_num_options + perf_counters.size()); context::pop_client(1); hsa_shut_down(); -} \ No newline at end of file +} diff --git a/tests/thread-trace/common.hpp b/tests/thread-trace/common.hpp index 10d4e5fc..efbd06a4 100644 --- a/tests/thread-trace/common.hpp +++ b/tests/thread-trace/common.hpp @@ -1,4 +1,27 @@ +// MIT License +// +// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// 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 +// AUTHORS 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 IN THE +// SOFTWARE. + #pragma once + #include #include @@ -100,4 +123,4 @@ callbacks_fini(); }; // namespace Callbacks -}; // namespace ATTTest \ No newline at end of file +}; // namespace ATTTest diff --git a/tests/thread-trace/main.cpp b/tests/thread-trace/main.cpp index 507fde9b..242bc433 100644 --- a/tests/thread-trace/main.cpp +++ b/tests/thread-trace/main.cpp @@ -92,4 +92,4 @@ main(int argc, char** argv) } return 0; -} \ No newline at end of file +} diff --git a/tests/thread-trace/trace_callbacks.cpp b/tests/thread-trace/trace_callbacks.cpp index 531e0add..97a319e3 100644 --- a/tests/thread-trace/trace_callbacks.cpp +++ b/tests/thread-trace/trace_callbacks.cpp @@ -238,4 +238,4 @@ callbacks_fini() } } // namespace Callbacks -} // namespace ATTTest \ No newline at end of file +} // namespace ATTTest