diff --git a/CHANGELOG.md b/CHANGELOG.md index 9b5b3f2caa..c65f2af7ff 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -104,6 +104,8 @@ Full documentation for ROCprofiler-SDK is available at [Click Here](source/docs/ - Replaced deprecated hipHostMalloc and hipHostFree functions with hipExtHostAlloc and hipFreeHost in when ROCm version is greater than or equal to 6.3 - Updated `rocprofv3` `--help` options. - Adding start and end timestamp columns to the counter collection csv output. +- Changed naming of agent profiling to device counting service (which more closely follows its name). To convert existing tool/user code to the new names, the following sed can be used: `find . -type f -exec sed -i 's/rocprofiler_agent_profile_callback_t/rocprofiler_device_counting_service_callback_t/g; s/rocprofiler_configure_agent_profile_counting_service/rocprofiler_configure_device_counting_service/g; s/agent_profile.h/device_counting_service.h/g; s/rocprofiler_sample_agent_profile_counting_service/rocprofiler_sample_device_counting_service/g' {} +` +- Changed naming of dispatch profiling service to dispatch counting service (which more closely follows its name). To convert existing tool/user code to the new names, the following sed can be used: `-type f -exec sed -i -e 's/dispatch_profile_counting_service/dispatch_counting_service/g' -e 's/dispatch_profile.h/dispatch_counting_service.h/g' -e 's/rocprofiler_profile_counting_dispatch_callback_t/rocprofiler_dispatch_counting_service_callback_t/g' -e 's/rocprofiler_profile_counting_dispatch_data_t/rocprofiler_dispatch_counting_service_data_t/g' -e 's/rocprofiler_profile_counting_dispatch_record_t/rocprofiler_dispatch_counting_service_record_t/g' {} +` ### Fixes diff --git a/samples/counter_collection/CMakeLists.txt b/samples/counter_collection/CMakeLists.txt index ea5797690f..d0906ec19b 100644 --- a/samples/counter_collection/CMakeLists.txt +++ b/samples/counter_collection/CMakeLists.txt @@ -107,28 +107,29 @@ set_tests_properties( "${counter-collection-functional-counter-env}" FAIL_REGULAR_EXPRESSION "${ROCPROFILER_DEFAULT_FAIL_REGEX}") -add_library(counter-collection-agent-profiling-client SHARED) -target_sources(counter-collection-agent-profiling-client PRIVATE agent_profiling.cpp - client.hpp) +add_library(counter-collection-device-profiling-client SHARED) +target_sources(counter-collection-device-profiling-client PRIVATE device_counting.cpp + client.hpp) target_link_libraries( - counter-collection-agent-profiling-client + counter-collection-device-profiling-client PUBLIC rocprofiler-sdk::samples-build-flags PRIVATE rocprofiler-sdk::rocprofiler-sdk rocprofiler-sdk::samples-common-library) -add_executable(counter-collection-agent-profiling) -target_sources(counter-collection-agent-profiling PRIVATE main.cpp) -target_link_libraries(counter-collection-agent-profiling - PRIVATE counter-collection-agent-profiling-client Threads::Threads) +add_executable(counter-collection-device-profiling) +target_sources(counter-collection-device-profiling PRIVATE main.cpp) +target_link_libraries(counter-collection-device-profiling + PRIVATE counter-collection-device-profiling-client Threads::Threads) -rocprofiler_samples_get_preload_env(PRELOAD_ENV counter-collection-agent-profiling-client) +rocprofiler_samples_get_preload_env(PRELOAD_ENV + counter-collection-device-profiling-client) set(counter-collection-functional-counter-env "${PRELOAD_ENV}" "${LIBRARY_PATH_ENV}") -add_test(NAME counter-collection-agent-profiling - COMMAND $) +add_test(NAME counter-collection-device-profiling + COMMAND $) set_tests_properties( - counter-collection-agent-profiling + counter-collection-device-profiling PROPERTIES TIMEOUT 120 LABELS "samples" ENVIRONMENT "${counter-collection-functional-counter-env}" FAIL_REGULAR_EXPRESSION "${ROCPROFILER_DEFAULT_FAIL_REGEX}") diff --git a/samples/counter_collection/callback_client.cpp b/samples/counter_collection/callback_client.cpp index 3e6a8ef24f..cf83091831 100644 --- a/samples/counter_collection/callback_client.cpp +++ b/samples/counter_collection/callback_client.cpp @@ -67,7 +67,7 @@ get_client_ctx() } void -record_callback(rocprofiler_profile_counting_dispatch_data_t dispatch_data, +record_callback(rocprofiler_dispatch_counting_service_data_t dispatch_data, rocprofiler_record_counter_t* record_data, size_t record_count, rocprofiler_user_data_t user_data, @@ -95,7 +95,7 @@ record_callback(rocprofiler_profile_counting_dispatch_data_t dispatch_data, * to collect the counter SQ_WAVES for all kernel dispatch packets. */ void -dispatch_callback(rocprofiler_profile_counting_dispatch_data_t dispatch_data, +dispatch_callback(rocprofiler_dispatch_counting_service_data_t dispatch_data, rocprofiler_profile_config_id_t* config, rocprofiler_user_data_t* /*user_data*/, void* /*callback_data_args*/) @@ -168,7 +168,7 @@ dispatch_callback(rocprofiler_profile_counting_dispatch_data_t dispatch_data, } // Create a colleciton profile for the counters - rocprofiler_profile_config_id_t profile; + rocprofiler_profile_config_id_t profile = {.handle = 0}; ROCPROFILER_CALL(rocprofiler_create_profile_config(dispatch_data.dispatch_info.agent_id, collect_counters.data(), collect_counters.size(), @@ -185,7 +185,7 @@ tool_init(rocprofiler_client_finalize_t, void* user_data) { ROCPROFILER_CALL(rocprofiler_create_context(&get_client_ctx()), "context creation failed"); - ROCPROFILER_CALL(rocprofiler_configure_callback_dispatch_profile_counting_service( + ROCPROFILER_CALL(rocprofiler_configure_callback_dispatch_counting_service( get_client_ctx(), dispatch_callback, nullptr, record_callback, user_data), "Could not setup counting service"); ROCPROFILER_CALL(rocprofiler_start_context(get_client_ctx()), "start context"); diff --git a/samples/counter_collection/client.cpp b/samples/counter_collection/client.cpp index aa4b5b38ad..3fd8a710ba 100644 --- a/samples/counter_collection/client.cpp +++ b/samples/counter_collection/client.cpp @@ -126,7 +126,7 @@ buffered_callback(rocprofiler_context_id_t, { // Print the returned counter data. auto* record = - static_cast(header->payload); + static_cast(header->payload); ss << "[Dispatch_Id: " << record->dispatch_info.dispatch_id << " Kernel_ID: " << record->dispatch_info.kernel_id << " Corr_Id: " << record->correlation_id.internal << ")]\n"; @@ -179,7 +179,7 @@ get_profile_cache() * to collect the counter SQ_WAVES for all kernel dispatch packets. */ void -dispatch_callback(rocprofiler_profile_counting_dispatch_data_t dispatch_data, +dispatch_callback(rocprofiler_dispatch_counting_service_data_t dispatch_data, rocprofiler_profile_config_id_t* config, rocprofiler_user_data_t* /*user_data*/, void* /*callback_data_args*/) @@ -255,7 +255,7 @@ build_profile_for_agent(rocprofiler_agent_id_t agent, } // Create and return the profile - rocprofiler_profile_config_id_t profile; + rocprofiler_profile_config_id_t profile = {.handle = 0}; ROCPROFILER_CALL(rocprofiler_create_profile_config( agent, collect_counters.data(), collect_counters.size(), &profile), "Could not construct profile cfg"); @@ -352,7 +352,7 @@ tool_init(rocprofiler_client_finalize_t, void* user_data) // counters to collect by returning a profile config id. In this example, we create the profile // configs above and store them in the map get_profile_cache() so we can look them up at // dispatch. - ROCPROFILER_CALL(rocprofiler_configure_buffered_dispatch_profile_counting_service( + ROCPROFILER_CALL(rocprofiler_configure_buffered_dispatch_counting_service( get_client_ctx(), get_buffer(), dispatch_callback, nullptr), "Could not setup buffered service"); diff --git a/samples/counter_collection/agent_profiling.cpp b/samples/counter_collection/device_counting.cpp similarity index 98% rename from samples/counter_collection/agent_profiling.cpp rename to samples/counter_collection/device_counting.cpp index b9a3435c72..a00b8b9e72 100644 --- a/samples/counter_collection/agent_profiling.cpp +++ b/samples/counter_collection/device_counting.cpp @@ -193,7 +193,7 @@ build_profile_for_agent(rocprofiler_agent_id_t agent) } } - rocprofiler_profile_config_id_t profile; + rocprofiler_profile_config_id_t profile = {.handle = 0}; ROCPROFILER_CALL(rocprofiler_create_profile_config( agent, collect_counters.data(), collect_counters.size(), &profile), "Could not construct profile cfg"); @@ -266,7 +266,7 @@ tool_init(rocprofiler_client_finalize_t, void* user_data) return 1; } - ROCPROFILER_CALL(rocprofiler_configure_agent_profile_counting_service( + ROCPROFILER_CALL(rocprofiler_configure_device_counting_service( get_client_ctx(), get_buffer(), agent_id, set_profile, nullptr), "Could not setup buffered service"); @@ -275,7 +275,7 @@ tool_init(rocprofiler_client_finalize_t, void* user_data) rocprofiler_start_context(get_client_ctx()); while(exit_toggle().load() == false) { - rocprofiler_sample_agent_profile_counting_service( + rocprofiler_sample_device_counting_service( get_client_ctx(), {.value = count}, ROCPROFILER_COUNTER_FLAG_NONE); count++; std::this_thread::sleep_for(std::chrono::milliseconds(50)); diff --git a/samples/counter_collection/print_functional_counters.cpp b/samples/counter_collection/print_functional_counters.cpp index af2b5e6fc8..665d781859 100644 --- a/samples/counter_collection/print_functional_counters.cpp +++ b/samples/counter_collection/print_functional_counters.cpp @@ -12,7 +12,7 @@ #include #include -#define PRINT_ONLY_FAILING true +#define PRINT_ONLY_FAILING false /** * Tests the collection of all counters on the agent the test is run on. @@ -222,7 +222,7 @@ get_agent_info() } void -dispatch_callback(rocprofiler_profile_counting_dispatch_data_t dispatch_data, +dispatch_callback(rocprofiler_dispatch_counting_service_data_t dispatch_data, rocprofiler_profile_config_id_t* config, rocprofiler_user_data_t* /*user_data*/, void* /*callback_data_args*/) @@ -304,7 +304,7 @@ dispatch_callback(rocprofiler_profile_counting_dispatch_data_t dispatch_data, } if(cap.remaining.empty()) return; - rocprofiler_profile_config_id_t profile; + rocprofiler_profile_config_id_t profile = {.handle = 0}; // Select the next counter to collect. if(rocprofiler_create_profile_config( @@ -312,6 +312,8 @@ dispatch_callback(rocprofiler_profile_counting_dispatch_data_t dispatch_data, ROCPROFILER_STATUS_SUCCESS) { *config = profile; + std::clog << "Attempting to read counter " + << cap.expected_counter_names.at(cap.remaining.back().handle) << "\n"; } cap.remaining.pop_back(); @@ -338,7 +340,7 @@ tool_init(rocprofiler_client_finalize_t, void*) ROCPROFILER_CALL(rocprofiler_assign_callback_thread(get_buffer(), client_thread), "failed to assign thread for buffer"); - ROCPROFILER_CALL(rocprofiler_configure_buffered_dispatch_profile_counting_service( + ROCPROFILER_CALL(rocprofiler_configure_buffered_dispatch_counting_service( get_client_ctx(), get_buffer(), dispatch_callback, nullptr), "Could not setup buffered service"); rocprofiler_start_context(get_client_ctx()); diff --git a/source/docs/api-reference/counter_collection_services.md b/source/docs/api-reference/counter_collection_services.md index 6d640aa053..3f6fe60700 100644 --- a/source/docs/api-reference/counter_collection_services.md +++ b/source/docs/api-reference/counter_collection_services.md @@ -51,14 +51,14 @@ After creating a context and buffer to store results, it is highly recommended ( // Setup the dispatch profile counting service. This service will trigger the dispatch_callback // when a kernel dispatch is enqueued into the HSA queue. The callback will specify what // counters to collect by returning a profile config id. - ROCPROFILER_CALL(rocprofiler_configure_buffered_dispatch_profile_counting_service( + ROCPROFILER_CALL(rocprofiler_configure_buffered_dispatch_counting_service( ctx, buff, dispatch_callback, nullptr), "Could not setup buffered service"); /* For Agent Profiling */ // set_profile is a callback that is use to select the profile to use when // the context is started. It is called at every rocprofiler_ctx_start() call. - ROCPROFILER_CALL(rocprofiler_configure_agent_profile_counting_service( + ROCPROFILER_CALL(rocprofiler_configure_device_counting_service( ctx, buff, agent_id, set_profile, nullptr), "Could not setup buffered service"); ``` @@ -162,13 +162,13 @@ When a kernel is dispatched, a dispatch callback is issued to the tool to allow ```CPP void -dispatch_callback(rocprofiler_profile_counting_dispatch_data_t dispatch_data, +dispatch_callback(rocprofiler_dispatch_counting_service_data_t dispatch_data, rocprofiler_profile_config_id_t* config, rocprofiler_user_data_t* user_data, void* /*callback_data_args*/) ``` -Dispatch data contains information about the dispatch that is being launched (such as its name) and config is where the tool can specify the profile (and in turn counters) to collect for the dispatch. If no profile is supplied, no counters are collected for this dispatch. User data contains user data supplied to rocprofiler_configure_buffered_dispatch_profile_counting_service. +Dispatch data contains information about the dispatch that is being launched (such as its name) and config is where the tool can specify the profile (and in turn counters) to collect for the dispatch. If no profile is supplied, no counters are collected for this dispatch. User data contains user data supplied to rocprofiler_configure_buffered_dispatch_counting_service. ### Agent Set Profile Callback @@ -197,7 +197,7 @@ Data from collected counter values is returned via a buffered callback. The buff { // Print the returned counter data. auto* record = - static_cast(header->payload); + static_cast(header->payload); ss << "[Dispatch_Id: " << record->dispatch_info.dispatch_id << " Kernel_ID: " << record->dispatch_info.kernel_id << " Corr_Id: " << record->correlation_id.internal << ")]\n"; diff --git a/source/include/rocprofiler-sdk/CMakeLists.txt b/source/include/rocprofiler-sdk/CMakeLists.txt index 68b57aab46..9fc64be4ce 100644 --- a/source/include/rocprofiler-sdk/CMakeLists.txt +++ b/source/include/rocprofiler-sdk/CMakeLists.txt @@ -11,14 +11,14 @@ set(ROCPROFILER_HEADER_FILES rocprofiler.h # secondary headers agent.h - agent_profile.h + device_counting_service.h buffer.h buffer_tracing.h callback_tracing.h context.h counters.h defines.h - dispatch_profile.h + dispatch_counting_service.h external_correlation.h fwd.h hip.h diff --git a/source/include/rocprofiler-sdk/cxx/serialization.hpp b/source/include/rocprofiler-sdk/cxx/serialization.hpp index e9f278076d..0a050e8f01 100644 --- a/source/include/rocprofiler-sdk/cxx/serialization.hpp +++ b/source/include/rocprofiler-sdk/cxx/serialization.hpp @@ -345,7 +345,7 @@ save(ArchiveT& ar, rocprofiler_callback_tracing_rccl_api_data_t data) template void -save(ArchiveT& ar, rocprofiler_profile_counting_dispatch_data_t data) +save(ArchiveT& ar, rocprofiler_dispatch_counting_service_data_t data) { ROCP_SDK_SAVE_DATA_FIELD(size); ROCP_SDK_SAVE_DATA_FIELD(correlation_id); @@ -356,7 +356,7 @@ save(ArchiveT& ar, rocprofiler_profile_counting_dispatch_data_t data) template void -save(ArchiveT& ar, rocprofiler_profile_counting_dispatch_record_t data) +save(ArchiveT& ar, rocprofiler_dispatch_counting_service_record_t data) { ROCP_SDK_SAVE_DATA_FIELD(size); ROCP_SDK_SAVE_DATA_FIELD(num_records); diff --git a/source/include/rocprofiler-sdk/agent_profile.h b/source/include/rocprofiler-sdk/device_counting_service.h similarity index 77% rename from source/include/rocprofiler-sdk/agent_profile.h rename to source/include/rocprofiler-sdk/device_counting_service.h index caad421aef..a0331be8f6 100644 --- a/source/include/rocprofiler-sdk/agent_profile.h +++ b/source/include/rocprofiler-sdk/device_counting_service.h @@ -26,7 +26,7 @@ #include /** - * @defgroup AGENT_PROFILE_COUNTING_SERVICE Agent Profile Counting Service + * @defgroup device_counting_service Agent Profile Counting Service * @brief needs brief description * * @{ @@ -59,23 +59,23 @@ typedef rocprofiler_status_t (*rocprofiler_agent_set_profile_callback_t)( * @param [in] agent_id agent id * @param [in] set_config Function to call to set the profile config (see * rocprofiler_agent_set_profile_callback_t) - * @param [in] user_data Data supplied to rocprofiler_configure_agent_profile_counting_service + * @param [in] user_data Data supplied to rocprofiler_configure_device_counting_service */ -typedef void (*rocprofiler_agent_profile_callback_t)( +typedef void (*rocprofiler_device_counting_service_callback_t)( rocprofiler_context_id_t context_id, rocprofiler_agent_id_t agent_id, rocprofiler_agent_set_profile_callback_t set_config, void* user_data); /** - * @brief Configure Profile Counting Service for agent. There may only be one agent profile - * configured per context and can be only one active context that is profiling a single agent - * at a time. Multiple agent contexts can be started at the same time if they are profiling + * @brief Configure Device Counting Service for agent. There may only be one counting service + * configured per agent in a context and can be only one active context that is profiling a single + * agent at a time. Multiple agent contexts can be started at the same time if they are profiling * different agents. * * @param [in] context_id context id * @param [in] buffer_id id of the buffer to use for the counting service. When - * rocprofiler_sample_agent_profile_counting_service is called, counter data will be written + * rocprofiler_sample_device_counting_service is called, counter data will be written * to this buffer. * @param [in] agent_id agent to configure profiling on. * @param [in] cb Callback called when the context is started for the tool to specify what @@ -87,16 +87,16 @@ typedef void (*rocprofiler_agent_profile_callback_t)( * @retval ::ROCPROFILER_STATUS_SUCCESS Returned if succesfully configured */ rocprofiler_status_t -rocprofiler_configure_agent_profile_counting_service(rocprofiler_context_id_t context_id, - rocprofiler_buffer_id_t buffer_id, - rocprofiler_agent_id_t agent_id, - rocprofiler_agent_profile_callback_t cb, - void* user_data) +rocprofiler_configure_device_counting_service(rocprofiler_context_id_t context_id, + rocprofiler_buffer_id_t buffer_id, + rocprofiler_agent_id_t agent_id, + rocprofiler_device_counting_service_callback_t cb, + void* user_data) ROCPROFILER_NONNULL(4) ROCPROFILER_API; /** * @brief Trigger a read of the counter data for the agent profile. The counter data will be - * written to the buffer specified in rocprofiler_configure_agent_profile_counting_service. + * written to the buffer specified in rocprofiler_configure_device_counting_service. * The data in rocprofiler_user_data_t will be written to the buffer along with the counter data. * flags can be used to specify if this call should be performed asynchronously (default is * synchronous). @@ -113,9 +113,9 @@ rocprofiler_configure_agent_profile_counting_service(rocprofiler_context_id_t co * @retval ::ROCPROFILER_STATUS_SUCCESS Returned if read request was successful. */ rocprofiler_status_t -rocprofiler_sample_agent_profile_counting_service(rocprofiler_context_id_t context_id, - rocprofiler_user_data_t user_data, - rocprofiler_counter_flag_t flags) ROCPROFILER_API; +rocprofiler_sample_device_counting_service(rocprofiler_context_id_t context_id, + rocprofiler_user_data_t user_data, + rocprofiler_counter_flag_t flags) ROCPROFILER_API; /** @} */ diff --git a/source/include/rocprofiler-sdk/dispatch_profile.h b/source/include/rocprofiler-sdk/dispatch_counting_service.h similarity index 83% rename from source/include/rocprofiler-sdk/dispatch_profile.h rename to source/include/rocprofiler-sdk/dispatch_counting_service.h index 296689149e..f881969ca1 100644 --- a/source/include/rocprofiler-sdk/dispatch_profile.h +++ b/source/include/rocprofiler-sdk/dispatch_counting_service.h @@ -31,7 +31,7 @@ ROCPROFILER_EXTERN_C_INIT /** - * @defgroup DISPATCH_PROFILE_COUNTING_SERVICE Dispatch Profile Counting Service + * @defgroup dispatch_counting_service Dispatch Profile Counting Service * @brief Per-dispatch hardware counter collection service * * @{ @@ -41,21 +41,21 @@ ROCPROFILER_EXTERN_C_INIT * @brief Kernel dispatch data for profile counting callbacks. * */ -typedef struct rocprofiler_profile_counting_dispatch_data_t +typedef struct rocprofiler_dispatch_counting_service_data_t { uint64_t size; ///< Size of this struct rocprofiler_correlation_id_t correlation_id; ///< Correlation ID for this dispatch rocprofiler_timestamp_t start_timestamp; ///< start time in nanoseconds rocprofiler_timestamp_t end_timestamp; ///< end time in nanoseconds rocprofiler_kernel_dispatch_info_t dispatch_info; ///< Dispatch info -} rocprofiler_profile_counting_dispatch_data_t; +} rocprofiler_dispatch_counting_service_data_t; /** * @brief ROCProfiler Profile Counting Counter Record Header Information * - * This is buffer equivalent of ::rocprofiler_profile_counting_dispatch_data_t + * This is buffer equivalent of ::rocprofiler_dispatch_counting_service_data_t */ -typedef struct rocprofiler_profile_counting_dispatch_record_t +typedef struct rocprofiler_dispatch_counting_service_record_t { uint64_t size; ///< Size of this struct uint64_t num_records; ///< number of ::rocprofiler_record_counter_t records @@ -63,7 +63,7 @@ typedef struct rocprofiler_profile_counting_dispatch_record_t rocprofiler_timestamp_t start_timestamp; ///< start time in nanoseconds rocprofiler_timestamp_t end_timestamp; ///< end time in nanoseconds rocprofiler_kernel_dispatch_info_t dispatch_info; ///< Contains the `dispatch_id` -} rocprofiler_profile_counting_dispatch_record_t; +} rocprofiler_dispatch_counting_service_record_t; /** * @brief Kernel Dispatch Callback. This is a callback that is invoked before the kernel @@ -72,13 +72,13 @@ typedef struct rocprofiler_profile_counting_dispatch_record_t * will be collected and emplaced in the buffer with @ref rocprofiler_buffer_id_t used when * setting up this callback. * - * @param [in] dispatch_data @see ::rocprofiler_profile_counting_dispatch_data_t + * @param [in] dispatch_data @see ::rocprofiler_dispatch_counting_service_data_t * @param [out] config Profile config detailing the counters to collect for this kernel * @param [out] user_data User data unique to this dispatch. Returned in record callback - * @param [in] callback_data_args Callback supplied via buffered_dispatch_profile_counting_service + * @param [in] callback_data_args Callback supplied via buffered_dispatch_counting_service */ -typedef void (*rocprofiler_profile_counting_dispatch_callback_t)( - rocprofiler_profile_counting_dispatch_data_t dispatch_data, +typedef void (*rocprofiler_dispatch_counting_service_callback_t)( + rocprofiler_dispatch_counting_service_data_t dispatch_data, rocprofiler_profile_config_id_t* config, rocprofiler_user_data_t* user_data, void* callback_data_args); @@ -86,17 +86,17 @@ typedef void (*rocprofiler_profile_counting_dispatch_callback_t)( /** * @brief Counting record callback. This is a callback is invoked when the kernel * execution is complete and contains the counter profile data requested in - * @ref rocprofiler_profile_counting_dispatch_callback_t. Only used with - * @ref rocprofiler_configure_callback_dispatch_profile_counting_service. + * @ref rocprofiler_dispatch_counting_service_callback_t. Only used with + * @ref rocprofiler_configure_callback_dispatch_counting_service. * - * @param [in] dispatch_data @see ::rocprofiler_profile_counting_dispatch_data_t + * @param [in] dispatch_data @see ::rocprofiler_dispatch_counting_service_data_t * @param [in] record_data Counter record data. * @param [in] record_count Number of counter records. * @param [in] user_data User data instance from dispatch callback - * @param [in] callback_data_args Callback supplied via buffered_dispatch_profile_counting_service + * @param [in] callback_data_args Callback supplied via buffered_dispatch_counting_service */ typedef void (*rocprofiler_profile_counting_record_callback_t)( - rocprofiler_profile_counting_dispatch_data_t dispatch_data, + rocprofiler_dispatch_counting_service_data_t dispatch_data, rocprofiler_record_counter_t* record_data, size_t record_count, rocprofiler_user_data_t user_data, @@ -112,14 +112,14 @@ typedef void (*rocprofiler_profile_counting_record_callback_t)( * NOTE: Interface is up for comment as to whether restrictions * on agent should be made here (limiting the CB based on agent) * or if the restriction should be performed by the tool in - * @ref rocprofiler_profile_counting_dispatch_callback_t (i.e. + * @ref rocprofiler_dispatch_counting_service_callback_t (i.e. * tool code checking the agent param to see if they want to profile * it). * * Interface is up for comment as to whether restrictions * on agent should be made here (limiting the CB based on agent) * or if the restriction should be performed by the tool in - * @ref rocprofiler_profile_counting_dispatch_callback_t (i.e. + * @ref rocprofiler_dispatch_counting_service_callback_t (i.e. * tool code checking the agent param to see if they want to profile * it). * @@ -130,10 +130,10 @@ typedef void (*rocprofiler_profile_counting_record_callback_t)( * @return ::rocprofiler_status_t */ rocprofiler_status_t -rocprofiler_configure_buffered_dispatch_profile_counting_service( +rocprofiler_configure_buffered_dispatch_counting_service( rocprofiler_context_id_t context_id, rocprofiler_buffer_id_t buffer_id, - rocprofiler_profile_counting_dispatch_callback_t callback, + rocprofiler_dispatch_counting_service_callback_t callback, void* callback_data_args) ROCPROFILER_API; /** @@ -149,9 +149,9 @@ rocprofiler_configure_buffered_dispatch_profile_counting_service( * @return ::rocprofiler_status_t */ rocprofiler_status_t -rocprofiler_configure_callback_dispatch_profile_counting_service( +rocprofiler_configure_callback_dispatch_counting_service( rocprofiler_context_id_t context_id, - rocprofiler_profile_counting_dispatch_callback_t dispatch_callback, + rocprofiler_dispatch_counting_service_callback_t dispatch_callback, void* dispatch_callback_args, rocprofiler_profile_counting_record_callback_t record_callback, void* record_callback_args) ROCPROFILER_API; diff --git a/source/include/rocprofiler-sdk/fwd.h b/source/include/rocprofiler-sdk/fwd.h index 0015493387..3c8c0ce826 100644 --- a/source/include/rocprofiler-sdk/fwd.h +++ b/source/include/rocprofiler-sdk/fwd.h @@ -390,13 +390,13 @@ typedef enum typedef enum { ROCPROFILER_COUNTER_RECORD_NONE = 0, - ROCPROFILER_COUNTER_RECORD_PROFILE_COUNTING_DISPATCH_HEADER, ///< ::rocprofiler_profile_counting_dispatch_record_t + ROCPROFILER_COUNTER_RECORD_PROFILE_COUNTING_DISPATCH_HEADER, ///< ::rocprofiler_dispatch_counting_service_record_t ROCPROFILER_COUNTER_RECORD_VALUE, ROCPROFILER_COUNTER_RECORD_LAST, /// @var ROCPROFILER_COUNTER_RECORD_KIND_DISPATCH_PROFILE_HEADER /// @brief Indicates the payload type is of type - /// ::rocprofiler_profile_counting_dispatch_record_t + /// ::rocprofiler_dispatch_counting_service_record_t } rocprofiler_counter_record_kind_t; /** @@ -694,9 +694,9 @@ typedef struct /// specific dispatch. /// /// This value can be mapped to a dispatch via the `dispatch_info` field (@see - /// ::rocprofiler_kernel_dispatch_info_t) of a ::rocprofiler_profile_counting_dispatch_data_t + /// ::rocprofiler_kernel_dispatch_info_t) of a ::rocprofiler_dispatch_counting_service_data_t /// instance (provided during callback for profile config) or a - /// ::rocprofiler_profile_counting_dispatch_record_t records (which will be insert into the + /// ::rocprofiler_dispatch_counting_service_record_t records (which will be insert into the /// buffer prior to the associated ::rocprofiler_record_counter_t records). } rocprofiler_record_counter_t; diff --git a/source/include/rocprofiler-sdk/rocprofiler.h b/source/include/rocprofiler-sdk/rocprofiler.h index bbf0e3ae8c..1681cef966 100644 --- a/source/include/rocprofiler-sdk/rocprofiler.h +++ b/source/include/rocprofiler-sdk/rocprofiler.h @@ -65,13 +65,13 @@ ROCPROFILER_EXTERN_C_FINI /** @} */ #include "rocprofiler-sdk/agent.h" -#include "rocprofiler-sdk/agent_profile.h" #include "rocprofiler-sdk/buffer.h" #include "rocprofiler-sdk/buffer_tracing.h" #include "rocprofiler-sdk/callback_tracing.h" #include "rocprofiler-sdk/context.h" #include "rocprofiler-sdk/counters.h" -#include "rocprofiler-sdk/dispatch_profile.h" +#include "rocprofiler-sdk/device_counting_service.h" +#include "rocprofiler-sdk/dispatch_counting_service.h" #include "rocprofiler-sdk/external_correlation.h" #include "rocprofiler-sdk/hip.h" #include "rocprofiler-sdk/hsa.h" diff --git a/source/lib/rocprofiler-sdk-tool/helper.hpp b/source/lib/rocprofiler-sdk-tool/helper.hpp index 2a001b436f..79f3441ba6 100644 --- a/source/lib/rocprofiler-sdk-tool/helper.hpp +++ b/source/lib/rocprofiler-sdk-tool/helper.hpp @@ -271,7 +271,7 @@ struct rocprofiler_tool_record_counter_t struct rocprofiler_tool_counter_collection_record_t { - rocprofiler_profile_counting_dispatch_data_t dispatch_data = {}; + rocprofiler_dispatch_counting_service_data_t dispatch_data = {}; std::array records = {}; uint64_t thread_id = 0; uint64_t arch_vgpr_count = 0; diff --git a/source/lib/rocprofiler-sdk-tool/tool.cpp b/source/lib/rocprofiler-sdk-tool/tool.cpp index f7dd3c4b38..acba07fa29 100644 --- a/source/lib/rocprofiler-sdk-tool/tool.cpp +++ b/source/lib/rocprofiler-sdk-tool/tool.cpp @@ -889,7 +889,7 @@ get_tool_agent(rocprofiler_agent_id_t id, const tool_agent_vec_t& data) // this function creates a rocprofiler profile config on the first entry auto -get_agent_profile(rocprofiler_agent_id_t agent_id) +get_device_counting_service(rocprofiler_agent_id_t agent_id) { static auto data = common::Synchronized{}; static const auto gpu_agents = get_gpu_agents(); @@ -980,7 +980,7 @@ get_agent_profile(rocprofiler_agent_id_t agent_id) } void -dispatch_callback(rocprofiler_profile_counting_dispatch_data_t dispatch_data, +dispatch_callback(rocprofiler_dispatch_counting_service_data_t dispatch_data, rocprofiler_profile_config_id_t* config, rocprofiler_user_data_t* user_data, void* /*callback_data_args*/) @@ -1004,7 +1004,7 @@ dispatch_callback(rocprofiler_profile_counting_dispatch_data_t dispatch_data, { return; } - else if(auto profile = get_agent_profile(agent_id)) + else if(auto profile = get_device_counting_service(agent_id)) { *config = *profile; user_data->value = common::get_tid(); @@ -1028,7 +1028,7 @@ get_counter_info_name(uint64_t record_id) } void -counter_record_callback(rocprofiler_profile_counting_dispatch_data_t dispatch_data, +counter_record_callback(rocprofiler_dispatch_counting_service_data_t dispatch_data, rocprofiler_record_counter_t* record_data, size_t record_count, rocprofiler_user_data_t user_data, @@ -1459,7 +1459,7 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data) if(tool::get_config().counter_collection) { ROCPROFILER_CALL( - rocprofiler_configure_callback_dispatch_profile_counting_service( + rocprofiler_configure_callback_dispatch_counting_service( get_client_ctx(), dispatch_callback, nullptr, counter_record_callback, nullptr), "Could not setup counting service"); } diff --git a/source/lib/rocprofiler-sdk/CMakeLists.txt b/source/lib/rocprofiler-sdk/CMakeLists.txt index aacd9281e8..d8d914bbd4 100644 --- a/source/lib/rocprofiler-sdk/CMakeLists.txt +++ b/source/lib/rocprofiler-sdk/CMakeLists.txt @@ -9,7 +9,7 @@ set(ROCPROFILER_LIB_SOURCES agent.cpp buffer.cpp buffer_tracing.cpp - agent_profile.cpp + device_counting_service.cpp callback_tracing.cpp context.cpp counters.cpp diff --git a/source/lib/rocprofiler-sdk/context/context.cpp b/source/lib/rocprofiler-sdk/context/context.cpp index 1a6d429fba..de310e88c8 100644 --- a/source/lib/rocprofiler-sdk/context/context.cpp +++ b/source/lib/rocprofiler-sdk/context/context.cpp @@ -323,7 +323,7 @@ start_context(rocprofiler_context_id_t context_id) if(cfg->counter_collection) rocprofiler::counters::start_context(cfg); if(cfg->agent_thread_trace) cfg->agent_thread_trace->start_context(); if(cfg->dispatch_thread_trace) cfg->dispatch_thread_trace->start_context(); - if(cfg->agent_counter_collection) status = rocprofiler::counters::start_agent_ctx(cfg); + if(cfg->device_counter_collection) status = rocprofiler::counters::start_agent_ctx(cfg); #if ROCPROFILER_SDK_HSA_PC_SAMPLING > 0 if(cfg->pc_sampler) status = rocprofiler::pc_sampling::start_service(cfg); #endif @@ -360,7 +360,7 @@ stop_context(rocprofiler_context_id_t idx) if(_expected->dispatch_thread_trace) _expected->dispatch_thread_trace->stop_context(); - if(_expected->agent_counter_collection) + if(_expected->device_counter_collection) { rocprofiler::counters::stop_agent_ctx(const_cast(_expected)); } diff --git a/source/lib/rocprofiler-sdk/context/context.hpp b/source/lib/rocprofiler-sdk/context/context.hpp index 168d468aa5..dd5cfb78fc 100644 --- a/source/lib/rocprofiler-sdk/context/context.hpp +++ b/source/lib/rocprofiler-sdk/context/context.hpp @@ -30,8 +30,8 @@ #include "lib/common/synchronized.hpp" #include "lib/rocprofiler-sdk/context/correlation_id.hpp" #include "lib/rocprofiler-sdk/context/domain.hpp" -#include "lib/rocprofiler-sdk/counters/agent_profiling.hpp" #include "lib/rocprofiler-sdk/counters/core.hpp" +#include "lib/rocprofiler-sdk/counters/device_counting.hpp" #include "lib/rocprofiler-sdk/external_correlation.hpp" #include "lib/rocprofiler-sdk/pc_sampling/types.hpp" #include "lib/rocprofiler-sdk/thread_trace/att_core.hpp" @@ -89,7 +89,7 @@ struct dispatch_counter_collection_service common::Synchronized enabled{false}; }; -struct agent_counter_collection_service +struct device_counting_service { std::vector agent_data; @@ -97,7 +97,8 @@ struct agent_counter_collection_service { DISABLED, LOCKED, - ENABLED + ENABLED, + EXIT }; std::atomic status{state::DISABLED}; @@ -124,9 +125,9 @@ struct context std::unique_ptr callback_tracer = {}; std::unique_ptr buffered_tracer = {}; // Only one of counter collection/agent counter collection can exists in the ctx. - std::unique_ptr counter_collection = {}; - std::unique_ptr agent_counter_collection = {}; - std::unique_ptr pc_sampler = {}; + std::unique_ptr counter_collection = {}; + std::unique_ptr device_counter_collection = {}; + std::unique_ptr pc_sampler = {}; std::unique_ptr dispatch_thread_trace = {}; std::unique_ptr agent_thread_trace = {}; diff --git a/source/lib/rocprofiler-sdk/counters/CMakeLists.txt b/source/lib/rocprofiler-sdk/counters/CMakeLists.txt index ff442d7db6..e38afffe9a 100644 --- a/source/lib/rocprofiler-sdk/counters/CMakeLists.txt +++ b/source/lib/rocprofiler-sdk/counters/CMakeLists.txt @@ -1,9 +1,9 @@ set(ROCPROFILER_LIB_COUNTERS_SOURCES metrics.cpp dimensions.cpp evaluate_ast.cpp core.cpp id_decode.cpp - dispatch_handlers.cpp controller.cpp agent_profiling.cpp) + dispatch_handlers.cpp controller.cpp device_counting.cpp) set(ROCPROFILER_LIB_COUNTERS_HEADERS metrics.hpp dimensions.hpp evaluate_ast.hpp core.hpp id_decode.hpp - dispatch_handlers.hpp controller.hpp agent_profiling.hpp) + dispatch_handlers.hpp controller.hpp device_counting.hpp) target_sources(rocprofiler-object-library PRIVATE ${ROCPROFILER_LIB_COUNTERS_SOURCES} ${ROCPROFILER_LIB_COUNTERS_HEADERS}) diff --git a/source/lib/rocprofiler-sdk/counters/controller.cpp b/source/lib/rocprofiler-sdk/counters/controller.cpp index 4ea09c33c0..60f59c7cab 100644 --- a/source/lib/rocprofiler-sdk/counters/controller.cpp +++ b/source/lib/rocprofiler-sdk/counters/controller.cpp @@ -23,7 +23,7 @@ #include "lib/rocprofiler-sdk/counters/controller.hpp" #include -#include +#include #include #include @@ -64,11 +64,11 @@ CounterController::destroy_profile(uint64_t id) } rocprofiler_status_t -CounterController::configure_agent_collection(rocprofiler_context_id_t context_id, - rocprofiler_buffer_id_t buffer_id, - rocprofiler_agent_id_t agent_id, - rocprofiler_agent_profile_callback_t cb, - void* user_data) +CounterController::configure_agent_collection(rocprofiler_context_id_t context_id, + rocprofiler_buffer_id_t buffer_id, + rocprofiler_agent_id_t agent_id, + rocprofiler_device_counting_service_callback_t cb, + void* user_data) { auto* ctx_p = rocprofiler::context::get_mutable_registered_context(context_id); if(!ctx_p) return ROCPROFILER_STATUS_ERROR_CONTEXT_INVALID; @@ -86,18 +86,18 @@ CounterController::configure_agent_collection(rocprofiler_context_id_t return ROCPROFILER_STATUS_ERROR_BUFFER_NOT_FOUND; } - if(!ctx.agent_counter_collection) + if(!ctx.device_counter_collection) { - ctx.agent_counter_collection = - std::make_unique(); + ctx.device_counter_collection = + std::make_unique(); } - ctx.agent_counter_collection->agent_data.emplace_back(); - ctx.agent_counter_collection->agent_data.back().callback_data = + ctx.device_counter_collection->agent_data.emplace_back(); + ctx.device_counter_collection->agent_data.back().callback_data = rocprofiler_user_data_t{.ptr = user_data}; - ctx.agent_counter_collection->agent_data.back().agent_id = agent_id; - ctx.agent_counter_collection->agent_data.back().cb = cb; - ctx.agent_counter_collection->agent_data.back().buffer = buffer_id; + ctx.device_counter_collection->agent_data.back().agent_id = agent_id; + ctx.device_counter_collection->agent_data.back().cb = cb; + ctx.device_counter_collection->agent_data.back().buffer = buffer_id; return ROCPROFILER_STATUS_SUCCESS; } @@ -110,7 +110,7 @@ rocprofiler_status_t CounterController::configure_dispatch( rocprofiler_context_id_t context_id, rocprofiler_buffer_id_t buffer, - rocprofiler_profile_counting_dispatch_callback_t callback, + rocprofiler_dispatch_counting_service_callback_t callback, void* callback_args, rocprofiler_profile_counting_record_callback_t record_callback, void* record_callback_args) @@ -120,7 +120,7 @@ CounterController::configure_dispatch( auto& ctx = *ctx_p; - if(ctx.agent_counter_collection) return ROCPROFILER_STATUS_ERROR_AGENT_DISPATCH_CONFLICT; + if(ctx.device_counter_collection) return ROCPROFILER_STATUS_ERROR_AGENT_DISPATCH_CONFLICT; // FIXME: Due to the clock gating issue, counter collection and PC sampling service // cannot coexist in the same context for now. diff --git a/source/lib/rocprofiler-sdk/counters/controller.hpp b/source/lib/rocprofiler-sdk/counters/controller.hpp index 488abaabe4..4acd904e81 100644 --- a/source/lib/rocprofiler-sdk/counters/controller.hpp +++ b/source/lib/rocprofiler-sdk/counters/controller.hpp @@ -29,7 +29,7 @@ #include "lib/rocprofiler-sdk/counters/metrics.hpp" #include -#include +#include #include #include @@ -81,17 +81,18 @@ public: static rocprofiler_status_t configure_dispatch( rocprofiler_context_id_t context_id, rocprofiler_buffer_id_t buffer, - rocprofiler_profile_counting_dispatch_callback_t callback, + rocprofiler_dispatch_counting_service_callback_t callback, void* callback_args, rocprofiler_profile_counting_record_callback_t record_callback, void* record_callback_args); std::shared_ptr get_profile_cfg(rocprofiler_profile_config_id_t id); - static rocprofiler_status_t configure_agent_collection(rocprofiler_context_id_t context_id, - rocprofiler_buffer_id_t buffer_id, - rocprofiler_agent_id_t agent_id, - rocprofiler_agent_profile_callback_t cb, - void* user_data); + static rocprofiler_status_t configure_agent_collection( + rocprofiler_context_id_t context_id, + rocprofiler_buffer_id_t buffer_id, + rocprofiler_agent_id_t agent_id, + rocprofiler_device_counting_service_callback_t cb, + void* user_data); private: rocprofiler::common::Synchronized>> diff --git a/source/lib/rocprofiler-sdk/counters/core.cpp b/source/lib/rocprofiler-sdk/counters/core.cpp index 065e05f37c..5df3629122 100644 --- a/source/lib/rocprofiler-sdk/counters/core.cpp +++ b/source/lib/rocprofiler-sdk/counters/core.cpp @@ -209,11 +209,11 @@ stop_context(const context::context* ctx) } rocprofiler_status_t -configure_agent_collection(rocprofiler_context_id_t context_id, - rocprofiler_buffer_id_t buffer_id, - rocprofiler_agent_id_t agent_id, - rocprofiler_agent_profile_callback_t cb, - void* user_data) +configure_agent_collection(rocprofiler_context_id_t context_id, + rocprofiler_buffer_id_t buffer_id, + rocprofiler_agent_id_t agent_id, + rocprofiler_device_counting_service_callback_t cb, + void* user_data) { return get_controller().configure_agent_collection( context_id, buffer_id, agent_id, cb, user_data); @@ -222,7 +222,7 @@ configure_agent_collection(rocprofiler_context_id_t context_id, rocprofiler_status_t configure_buffered_dispatch(rocprofiler_context_id_t context_id, rocprofiler_buffer_id_t buffer, - rocprofiler_profile_counting_dispatch_callback_t callback, + rocprofiler_dispatch_counting_service_callback_t callback, void* callback_args) { CHECK_NE(buffer.handle, 0); @@ -232,7 +232,7 @@ configure_buffered_dispatch(rocprofiler_context_id_t con rocprofiler_status_t configure_callback_dispatch(rocprofiler_context_id_t context_id, - rocprofiler_profile_counting_dispatch_callback_t callback, + rocprofiler_dispatch_counting_service_callback_t callback, void* callback_data_args, rocprofiler_profile_counting_record_callback_t record_callback, void* record_callback_args) diff --git a/source/lib/rocprofiler-sdk/counters/core.hpp b/source/lib/rocprofiler-sdk/counters/core.hpp index 04ead96342..85a88fd6b0 100644 --- a/source/lib/rocprofiler-sdk/counters/core.hpp +++ b/source/lib/rocprofiler-sdk/counters/core.hpp @@ -23,7 +23,7 @@ #pragma once #include -#include +#include #include #include @@ -45,7 +45,7 @@ namespace counters struct counter_callback_info { // User callback - rocprofiler_profile_counting_dispatch_callback_t user_cb{nullptr}; + rocprofiler_dispatch_counting_service_callback_t user_cb{nullptr}; // User id void* callback_args{nullptr}; // Link to the context this is associated with @@ -81,22 +81,22 @@ destroy_counter_profile(uint64_t id); rocprofiler_status_t configure_buffered_dispatch(rocprofiler_context_id_t context_id, rocprofiler_buffer_id_t buffer, - rocprofiler_profile_counting_dispatch_callback_t callback, + rocprofiler_dispatch_counting_service_callback_t callback, void* callback_args); rocprofiler_status_t configure_callback_dispatch(rocprofiler_context_id_t context_id, - rocprofiler_profile_counting_dispatch_callback_t callback, + rocprofiler_dispatch_counting_service_callback_t callback, void* callback_data_args, rocprofiler_profile_counting_record_callback_t record_callback, void* record_callback_args); rocprofiler_status_t -configure_agent_collection(rocprofiler_context_id_t context_id, - rocprofiler_buffer_id_t buffer_id, - rocprofiler_agent_id_t agent_id, - rocprofiler_agent_profile_callback_t cb, - void* user_data); +configure_agent_collection(rocprofiler_context_id_t context_id, + rocprofiler_buffer_id_t buffer_id, + rocprofiler_agent_id_t agent_id, + rocprofiler_device_counting_service_callback_t cb, + void* user_data); void start_context(const context::context*); diff --git a/source/lib/rocprofiler-sdk/counters/agent_profiling.cpp b/source/lib/rocprofiler-sdk/counters/device_counting.cpp similarity index 86% rename from source/lib/rocprofiler-sdk/counters/agent_profiling.cpp rename to source/lib/rocprofiler-sdk/counters/device_counting.cpp index 2a4c9a0cd2..06e5b8a2af 100644 --- a/source/lib/rocprofiler-sdk/counters/agent_profiling.cpp +++ b/source/lib/rocprofiler-sdk/counters/device_counting.cpp @@ -20,7 +20,7 @@ // OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE // SOFTWARE. -#include "lib/rocprofiler-sdk/counters/agent_profiling.hpp" +#include "lib/rocprofiler-sdk/counters/device_counting.hpp" #include "lib/common/logging.hpp" #include "lib/rocprofiler-sdk/buffer.hpp" #include "lib/rocprofiler-sdk/context/context.hpp" @@ -138,6 +138,13 @@ agent_async_handler(hsa_signal_value_t /*signal_v*/, void* data) return false; } + if(decoded_pkt.empty()) + { + // reset the signal to allow another sample to start + hsa::get_core_table()->hsa_signal_store_relaxed_fn(callback_data.completion, 1); + return true; + } + // Write out the AQL data to the buffer for(auto& ast : prof_config->asts) { @@ -249,13 +256,13 @@ read_agent_ctx(const context::context* ctx, rocprofiler_counter_flag_t flags) { rocprofiler_status_t status = ROCPROFILER_STATUS_SUCCESS; - if(!ctx->agent_counter_collection) + if(!ctx->device_counter_collection) { ROCP_ERROR << fmt::format("Context {} has no agent counter collection", ctx->context_idx); return ROCPROFILER_STATUS_ERROR_CONTEXT_INVALID; } - auto& agent_ctx = *ctx->agent_counter_collection; + auto& agent_ctx = *ctx->device_counter_collection; // If we have not initiualized HSA yet, nothing to read, return; if(hsa_inited().load() == false) @@ -264,9 +271,9 @@ read_agent_ctx(const context::context* ctx, } // Set the state to LOCKED to prevent other calls to start/stop/read. - auto expected = rocprofiler::context::agent_counter_collection_service::state::ENABLED; + auto expected = rocprofiler::context::device_counting_service::state::ENABLED; if(!agent_ctx.status.compare_exchange_strong( - expected, rocprofiler::context::agent_counter_collection_service::state::LOCKED)) + expected, rocprofiler::context::device_counting_service::state::LOCKED)) { return ROCPROFILER_STATUS_ERROR_CONTEXT_ERROR; } @@ -337,8 +344,7 @@ read_agent_ctx(const context::context* ctx, } } - agent_ctx.status.exchange( - rocprofiler::context::agent_counter_collection_service::state::ENABLED); + agent_ctx.status.exchange(rocprofiler::context::device_counting_service::state::ENABLED); return status; } @@ -357,12 +363,12 @@ rocprofiler_status_t start_agent_ctx(const context::context* ctx) { auto status = ROCPROFILER_STATUS_SUCCESS; - if(!ctx->agent_counter_collection) + if(!ctx->device_counter_collection) { return status; } - auto& agent_ctx = *ctx->agent_counter_collection; + auto& agent_ctx = *ctx->device_counter_collection; if(hsa_inited().load() == false) { @@ -370,9 +376,9 @@ start_agent_ctx(const context::context* ctx) } // Set the state to LOCKED to prevent other calls to start/stop/read. - auto expected = rocprofiler::context::agent_counter_collection_service::state::DISABLED; + auto expected = rocprofiler::context::device_counting_service::state::DISABLED; if(!agent_ctx.status.compare_exchange_strong( - expected, rocprofiler::context::agent_counter_collection_service::state::LOCKED)) + expected, rocprofiler::context::device_counting_service::state::LOCKED)) { return ROCPROFILER_STATUS_ERROR_SERVICE_ALREADY_CONFIGURED; } @@ -408,19 +414,19 @@ start_agent_ctx(const context::context* ctx) auto config = rocprofiler::counters::get_profile_config(config_id); if(!config) return ROCPROFILER_STATUS_ERROR_PROFILE_NOT_FOUND; - if(!cb_ctx->agent_counter_collection) + if(!cb_ctx->device_counter_collection) { return ROCPROFILER_STATUS_ERROR_CONTEXT_INVALID; } // Only allow profiles to be set in the locked state - if(cb_ctx->agent_counter_collection->status.load() != - rocprofiler::context::agent_counter_collection_service::state::LOCKED) + if(cb_ctx->device_counter_collection->status.load() != + rocprofiler::context::device_counting_service::state::LOCKED) { return ROCPROFILER_STATUS_ERROR_CONFIGURATION_LOCKED; } - for(auto& agent_data : cb_ctx->agent_counter_collection->agent_data) + for(auto& agent_data : cb_ctx->device_counter_collection->agent_data) { // Find the agent that this profile is for and set it. if(agent_data.agent_id.handle == config->agent->id.handle) @@ -474,8 +480,7 @@ start_agent_ctx(const context::context* ctx) HSA_WAIT_STATE_ACTIVE); } - agent_ctx.status.exchange( - rocprofiler::context::agent_counter_collection_service::state::ENABLED); + agent_ctx.status.exchange(rocprofiler::context::device_counting_service::state::ENABLED); return status; } @@ -490,21 +495,21 @@ rocprofiler_status_t stop_agent_ctx(const context::context* ctx) { auto status = ROCPROFILER_STATUS_SUCCESS; - if(!ctx->agent_counter_collection) + if(!ctx->device_counter_collection) { return status; } - auto& agent_ctx = *ctx->agent_counter_collection; + auto& agent_ctx = *ctx->device_counter_collection; if(hsa_inited().load() == false) { return ROCPROFILER_STATUS_SUCCESS; } - auto expected = rocprofiler::context::agent_counter_collection_service::state::ENABLED; + auto expected = rocprofiler::context::device_counting_service::state::ENABLED; if(!agent_ctx.status.compare_exchange_strong( - expected, rocprofiler::context::agent_counter_collection_service::state::LOCKED)) + expected, rocprofiler::context::device_counting_service::state::LOCKED)) { // Status is already stopped or being enabled elsewhere. return ROCPROFILER_STATUS_SUCCESS; @@ -531,21 +536,48 @@ stop_agent_ctx(const context::context* ctx) HSA_WAIT_STATE_ACTIVE); } - agent_ctx.status.exchange( - rocprofiler::context::agent_counter_collection_service::state::DISABLED); + agent_ctx.status.exchange(rocprofiler::context::device_counting_service::state::DISABLED); return status; } +// Stop all contexts and prevent any further requests to start/stop/read. +// Waits until any current operation is complete before exiting. +rocprofiler_status_t +device_counting_service_finalize() +{ + for(auto& ctx : context::get_registered_contexts()) + { + std::vector expected = { + rocprofiler::context::device_counting_service::state::DISABLED, + rocprofiler::context::device_counting_service::state::ENABLED, + rocprofiler::context::device_counting_service::state::EXIT}; + if(!ctx->device_counter_collection) continue; + while(!ctx->device_counter_collection->status.compare_exchange_strong( + expected[0], rocprofiler::context::device_counting_service::state::EXIT) && + !ctx->device_counter_collection->status.compare_exchange_strong( + expected[1], rocprofiler::context::device_counting_service::state::EXIT) && + !ctx->device_counter_collection->status.compare_exchange_strong( + expected[2], rocprofiler::context::device_counting_service::state::EXIT)) + { + // Note: Compare Exchange can modify expected even if the exchange fails + expected = {rocprofiler::context::device_counting_service::state::DISABLED, + rocprofiler::context::device_counting_service::state::ENABLED, + rocprofiler::context::device_counting_service::state::EXIT}; + }; + } + return ROCPROFILER_STATUS_SUCCESS; +} + // If we have ctx's that were started before HSA was initialized, we need to // actually start those contexts now that we have an HSA instance. rocprofiler_status_t -agent_profile_hsa_registration() +device_counting_service_hsa_registration() { hsa_inited().store(true); for(auto& ctx : context::get_active_contexts()) { - if(!ctx->agent_counter_collection) continue; + if(!ctx->device_counter_collection) continue; start_agent_ctx(ctx); } diff --git a/source/lib/rocprofiler-sdk/counters/agent_profiling.hpp b/source/lib/rocprofiler-sdk/counters/device_counting.hpp similarity index 93% rename from source/lib/rocprofiler-sdk/counters/agent_profiling.hpp rename to source/lib/rocprofiler-sdk/counters/device_counting.hpp index da520d76c4..8436bbc291 100644 --- a/source/lib/rocprofiler-sdk/counters/agent_profiling.hpp +++ b/source/lib/rocprofiler-sdk/counters/device_counting.hpp @@ -57,7 +57,7 @@ struct agent_callback_data std::shared_ptr profile = {}; rocprofiler_agent_id_t agent_id = {.handle = 0}; - rocprofiler_agent_profile_callback_t cb = nullptr; + rocprofiler_device_counting_service_callback_t cb = nullptr; rocprofiler_buffer_id_t buffer = {.handle = 0}; bool set_profile = false; @@ -78,11 +78,16 @@ struct agent_callback_data ~agent_callback_data(); }; +// Stop all contexts and prevent any further requests to start/stop/read. +// Waits until any current operation is complete before exiting. +rocprofiler_status_t +device_counting_service_finalize(); + // If we have contexts that are started before HSA init. This // function will start those contexts. Should only be called // as part of the HSA init process in rocprofiler. rocprofiler_status_t -agent_profile_hsa_registration(); +device_counting_service_hsa_registration(); // Send the AQL start packet to a queue on the agent to start // collecting counter data. This function is synchronous and will diff --git a/source/lib/rocprofiler-sdk/counters/dispatch_handlers.cpp b/source/lib/rocprofiler-sdk/counters/dispatch_handlers.cpp index 2031dc28e0..d59fdf4df4 100644 --- a/source/lib/rocprofiler-sdk/counters/dispatch_handlers.cpp +++ b/source/lib/rocprofiler-sdk/counters/dispatch_handlers.cpp @@ -108,7 +108,7 @@ queue_cb(const context::context* ctx, auto req_profile = rocprofiler_profile_config_id_t{.handle = 0}; auto dispatch_data = - common::init_public_api_struct(rocprofiler_profile_counting_dispatch_data_t{}); + common::init_public_api_struct(rocprofiler_dispatch_counting_service_data_t{}); dispatch_data.correlation_id = _corr_id_v; { @@ -248,7 +248,7 @@ completed_cb(const context::context* ctx, if(buf) { auto _header = - common::init_public_api_struct(rocprofiler_profile_counting_dispatch_record_t{}); + common::init_public_api_struct(rocprofiler_dispatch_counting_service_record_t{}); _header.num_records = out.size(); _header.correlation_id = _corr_id_v; if(dispatch_time.status == HSA_STATUS_SUCCESS) @@ -270,7 +270,7 @@ completed_cb(const context::context* ctx, CHECK(info->record_callback); auto dispatch_data = - common::init_public_api_struct(rocprofiler_profile_counting_dispatch_data_t{}); + common::init_public_api_struct(rocprofiler_dispatch_counting_service_data_t{}); dispatch_data.dispatch_info = session.callback_record.dispatch_info; dispatch_data.correlation_id = _corr_id_v; diff --git a/source/lib/rocprofiler-sdk/counters/evaluate_ast.cpp b/source/lib/rocprofiler-sdk/counters/evaluate_ast.cpp index cf68cb420a..96de015424 100644 --- a/source/lib/rocprofiler-sdk/counters/evaluate_ast.cpp +++ b/source/lib/rocprofiler-sdk/counters/evaluate_ast.cpp @@ -509,7 +509,11 @@ EvaluateAST::read_pkt(const aql::CounterPacketConstruct* pkt_gen, hsa::AQLPacket return HSA_STATUS_SUCCESS; }, &aql_data); - CHECK(status == HSA_STATUS_SUCCESS); + + if(status != HSA_STATUS_SUCCESS) + { + ROCP_ERROR << "AqlProfile could not decode packet"; + } return ret; } diff --git a/source/lib/rocprofiler-sdk/counters/tests/CMakeLists.txt b/source/lib/rocprofiler-sdk/counters/tests/CMakeLists.txt index e012dba287..4d476bffcc 100644 --- a/source/lib/rocprofiler-sdk/counters/tests/CMakeLists.txt +++ b/source/lib/rocprofiler-sdk/counters/tests/CMakeLists.txt @@ -56,8 +56,8 @@ target_link_libraries( set(ROCPROFILER_LIB_COUNTER_TEST_SOURCES metrics_test.cpp evaluate_ast_test.cpp dimension.cpp init_order.cpp core.cpp - code_object_loader.cpp agent_profiling.cpp) -set(ROCPROFILER_LIB_COUNTER_TEST_HEADERS code_object_loader.hpp agent_profiling.hpp) + code_object_loader.cpp device_counting.cpp) +set(ROCPROFILER_LIB_COUNTER_TEST_HEADERS code_object_loader.hpp device_counting.hpp) add_executable(counter-test) diff --git a/source/lib/rocprofiler-sdk/counters/tests/core.cpp b/source/lib/rocprofiler-sdk/counters/tests/core.cpp index 473baa9b45..83ea2f3eaa 100644 --- a/source/lib/rocprofiler-sdk/counters/tests/core.cpp +++ b/source/lib/rocprofiler-sdk/counters/tests/core.cpp @@ -33,7 +33,7 @@ #include "lib/rocprofiler-sdk/kernel_dispatch/profiling_time.hpp" #include "lib/rocprofiler-sdk/registration.hpp" -#include +#include #include #include #include @@ -165,7 +165,7 @@ buffered_callback(rocprofiler_context_id_t, } void -null_dispatch_callback(rocprofiler_profile_counting_dispatch_data_t, +null_dispatch_callback(rocprofiler_dispatch_counting_service_data_t, rocprofiler_profile_config_id_t*, rocprofiler_user_data_t*, void*) @@ -181,7 +181,7 @@ null_buffered_callback(rocprofiler_context_id_t, {} void -null_record_callback(rocprofiler_profile_counting_dispatch_data_t, +null_record_callback(rocprofiler_dispatch_counting_service_data_t, rocprofiler_record_counter_t*, size_t, rocprofiler_user_data_t, @@ -207,7 +207,7 @@ TEST(core, check_packet_generation) /** * Check profile construction */ - rocprofiler_profile_config_id_t cfg_id = {}; + rocprofiler_profile_config_id_t cfg_id = {.handle = 0}; rocprofiler_counter_id_t id = {.handle = metric.id()}; ROCP_ERROR << fmt::format("Generating packet for {}", metric); ROCPROFILER_CALL( @@ -301,7 +301,7 @@ namespace struct expected_dispatch { // To pass back - rocprofiler_profile_config_id_t id = {}; + rocprofiler_profile_config_id_t id = {.handle = 0}; rocprofiler_queue_id_t queue_id = {.handle = 0}; rocprofiler_agent_id_t agent_id = {.handle = 0}; uint64_t kernel_id = 0; @@ -313,7 +313,7 @@ struct expected_dispatch }; void -user_dispatch_cb(rocprofiler_profile_counting_dispatch_data_t dispatch_data, +user_dispatch_cb(rocprofiler_dispatch_counting_service_data_t dispatch_data, rocprofiler_profile_config_id_t* config, rocprofiler_user_data_t* user_data, void* callback_data_args) @@ -326,7 +326,7 @@ user_dispatch_cb(rocprofiler_profile_counting_dispatch_data_t dispatch_data, auto kernel_id = dispatch_data.dispatch_info.kernel_id; auto dispatch_id = dispatch_data.dispatch_info.dispatch_id; - EXPECT_EQ(sizeof(rocprofiler_profile_counting_dispatch_data_t), dispatch_data.size); + EXPECT_EQ(sizeof(rocprofiler_dispatch_counting_service_data_t), dispatch_data.size); EXPECT_EQ(expected.kernel_id, kernel_id); EXPECT_EQ(expected.dispatch_id, dispatch_id); EXPECT_EQ(expected.agent_id, agent_id); @@ -533,7 +533,7 @@ TEST(core, start_stop_buffered_ctx) &opt_buff_id), "Could not create buffer"); - ROCPROFILER_CALL(rocprofiler_configure_buffered_dispatch_profile_counting_service( + ROCPROFILER_CALL(rocprofiler_configure_buffered_dispatch_counting_service( get_client_ctx(), opt_buff_id, null_dispatch_callback, (void*) 0x12345), "Could not setup buffered service"); ROCPROFILER_CALL(rocprofiler_start_context(get_client_ctx()), "start context"); @@ -595,11 +595,11 @@ TEST(core, start_stop_callback_ctx) ROCPROFILER_CALL(rocprofiler_create_context(&get_client_ctx()), "context creation failed"); ROCPROFILER_CALL( - rocprofiler_configure_callback_dispatch_profile_counting_service(get_client_ctx(), - null_dispatch_callback, - (void*) 0x12345, - null_record_callback, - (void*) 0x54321), + rocprofiler_configure_callback_dispatch_counting_service(get_client_ctx(), + null_dispatch_callback, + (void*) 0x12345, + null_record_callback, + (void*) 0x54321), "Could not setup counting service"); ROCPROFILER_CALL(rocprofiler_start_context(get_client_ctx()), "start context"); diff --git a/source/lib/rocprofiler-sdk/counters/tests/agent_profiling.cpp b/source/lib/rocprofiler-sdk/counters/tests/device_counting.cpp similarity index 95% rename from source/lib/rocprofiler-sdk/counters/tests/agent_profiling.cpp rename to source/lib/rocprofiler-sdk/counters/tests/device_counting.cpp index 8af74cfa09..b3ec26ef67 100644 --- a/source/lib/rocprofiler-sdk/counters/tests/agent_profiling.cpp +++ b/source/lib/rocprofiler-sdk/counters/tests/device_counting.cpp @@ -33,7 +33,7 @@ #include "lib/rocprofiler-sdk/registration.hpp" #include -#include +#include #include #include #include @@ -241,10 +241,10 @@ submitPacket(hsa_queue_t* queue, const void* packet) } // namespace -class agent_profile_test : public ::testing::Test +class device_counting_service_test : public ::testing::Test { protected: - agent_profile_test() {} + device_counting_service_test() {} static void test_run(rocprofiler_counter_flag_t flags = ROCPROFILER_COUNTER_FLAG_NONE, const std::unordered_set& test_metrics = {}, @@ -256,7 +256,7 @@ protected: context::push_client(1); test_init(); // rocprofiler_debugger_block(); - counters::agent_profile_hsa_registration(); + counters::device_counting_service_hsa_registration(); std::string kernel_name = "null_kernel"; @@ -336,14 +336,14 @@ protected: /** * Check profile construction */ - rocprofiler_profile_config_id_t cfg_id = {}; + rocprofiler_profile_config_id_t cfg_id = {.handle = 0}; rocprofiler_counter_id_t id = {.handle = metric.id()}; ROCPROFILER_CALL( rocprofiler_create_profile_config(agent.get_rocp_agent()->id, &id, 1, &cfg_id), "Unable to create profile"); ROCPROFILER_CALL( - rocprofiler_configure_agent_profile_counting_service( + rocprofiler_configure_device_counting_service( ctx, opt_buff_id, agent.get_rocp_agent()->id, @@ -367,7 +367,7 @@ protected: // construction This is a workaround for the test environment since we create // contexts after AgentCache constructed. agent::get_agent_cache(agent.get_rocp_agent()) - ->init_agent_profile_queue(get_api_table(), get_ext_table()); + ->init_device_counting_service_queue(get_api_table(), get_ext_table()); hsa_signal_store_screlease(completion_signal, 1); hsa_signal_store_screlease(found_data, 0); @@ -402,9 +402,9 @@ protected: HSA_WAIT_STATE_BLOCKED); // Sample the counting service. - ROCPROFILER_CALL(rocprofiler_sample_agent_profile_counting_service( - ctx, {.value = track_metric}, flags), - "Could not sample"); + ROCPROFILER_CALL( + rocprofiler_sample_device_counting_service(ctx, {.value = track_metric}, flags), + "Could not sample"); ROCPROFILER_CALL(rocprofiler_stop_context(ctx), "Could not stop context"); rocprofiler_flush_buffer(opt_buff_id); @@ -427,9 +427,9 @@ protected: } }; -TEST_F(agent_profile_test, sync_counters) { test_run(); } -TEST_F(agent_profile_test, async_counters) { test_run(ROCPROFILER_COUNTER_FLAG_ASYNC); } -TEST_F(agent_profile_test, sync_grbm_verify) +TEST_F(device_counting_service_test, sync_counters) { test_run(); } +TEST_F(device_counting_service_test, async_counters) { test_run(ROCPROFILER_COUNTER_FLAG_ASYNC); } +TEST_F(device_counting_service_test, sync_grbm_verify) { test_run(ROCPROFILER_COUNTER_FLAG_NONE, {"GRBM_COUNT"}, 50000); ROCP_ERROR << global_recs().size(); @@ -445,7 +445,7 @@ TEST_F(agent_profile_test, sync_grbm_verify) } } -TEST_F(agent_profile_test, sync_gpu_util_verify) +TEST_F(device_counting_service_test, sync_gpu_util_verify) { test_run(ROCPROFILER_COUNTER_FLAG_NONE, {"GPU_UTIL"}, 50000); ROCP_ERROR << global_recs().size(); @@ -461,7 +461,7 @@ TEST_F(agent_profile_test, sync_gpu_util_verify) } } -TEST_F(agent_profile_test, sync_sq_waves_verify) +TEST_F(device_counting_service_test, sync_sq_waves_verify) { test_run(ROCPROFILER_COUNTER_FLAG_NONE, {"SQ_WAVES_sum"}, 50000); ROCP_ERROR << global_recs().size(); diff --git a/source/lib/rocprofiler-sdk/counters/tests/agent_profiling.hpp b/source/lib/rocprofiler-sdk/counters/tests/device_counting.hpp similarity index 100% rename from source/lib/rocprofiler-sdk/counters/tests/agent_profiling.hpp rename to source/lib/rocprofiler-sdk/counters/tests/device_counting.hpp diff --git a/source/lib/rocprofiler-sdk/counters/tests/init_order.cpp b/source/lib/rocprofiler-sdk/counters/tests/init_order.cpp index 06f297a219..98796f24db 100644 --- a/source/lib/rocprofiler-sdk/counters/tests/init_order.cpp +++ b/source/lib/rocprofiler-sdk/counters/tests/init_order.cpp @@ -122,7 +122,7 @@ buffered_callback(rocprofiler_context_id_t, {} void -dispatch_callback(rocprofiler_profile_counting_dispatch_data_t, +dispatch_callback(rocprofiler_dispatch_counting_service_data_t, rocprofiler_profile_config_id_t*, rocprofiler_user_data_t*, void*) @@ -163,7 +163,7 @@ TEST(counters_init_order, metric_map_order) nullptr, &get_buffer()), "buffer creation failed"); - ROCPROFILER_CALL(rocprofiler_configure_buffered_dispatch_profile_counting_service( + ROCPROFILER_CALL(rocprofiler_configure_buffered_dispatch_counting_service( get_client_ctx(), get_buffer(), dispatch_callback, nullptr), "Could not setup buffered service"); rocprofiler::registration::set_init_status(1); diff --git a/source/lib/rocprofiler-sdk/agent_profile.cpp b/source/lib/rocprofiler-sdk/device_counting_service.cpp similarity index 67% rename from source/lib/rocprofiler-sdk/agent_profile.cpp rename to source/lib/rocprofiler-sdk/device_counting_service.cpp index 44089eb634..449a93fda4 100644 --- a/source/lib/rocprofiler-sdk/agent_profile.cpp +++ b/source/lib/rocprofiler-sdk/device_counting_service.cpp @@ -23,26 +23,26 @@ #include #include "lib/rocprofiler-sdk/context/context.hpp" -#include "lib/rocprofiler-sdk/counters/agent_profiling.hpp" #include "lib/rocprofiler-sdk/counters/core.hpp" +#include "lib/rocprofiler-sdk/counters/device_counting.hpp" #include "rocprofiler-sdk/fwd.h" extern "C" { rocprofiler_status_t -rocprofiler_configure_agent_profile_counting_service(rocprofiler_context_id_t context_id, - rocprofiler_buffer_id_t buffer_id, - rocprofiler_agent_id_t agent_id, - rocprofiler_agent_profile_callback_t cb, - void* user_data) +rocprofiler_configure_device_counting_service(rocprofiler_context_id_t context_id, + rocprofiler_buffer_id_t buffer_id, + rocprofiler_agent_id_t agent_id, + rocprofiler_device_counting_service_callback_t cb, + void* user_data) { return rocprofiler::counters::configure_agent_collection( context_id, buffer_id, agent_id, cb, user_data); } rocprofiler_status_t -rocprofiler_sample_agent_profile_counting_service(rocprofiler_context_id_t context_id, - rocprofiler_user_data_t user_data, - rocprofiler_counter_flag_t flags) +rocprofiler_sample_device_counting_service(rocprofiler_context_id_t context_id, + rocprofiler_user_data_t user_data, + rocprofiler_counter_flag_t flags) { return rocprofiler::counters::read_agent_ctx( rocprofiler::context::get_registered_context(context_id), user_data, flags); diff --git a/source/lib/rocprofiler-sdk/dispatch_profile.cpp b/source/lib/rocprofiler-sdk/dispatch_profile.cpp index 8bbf8f171e..8e8d822eab 100644 --- a/source/lib/rocprofiler-sdk/dispatch_profile.cpp +++ b/source/lib/rocprofiler-sdk/dispatch_profile.cpp @@ -39,10 +39,10 @@ extern "C" { * @return ::rocprofiler_status_t */ rocprofiler_status_t -rocprofiler_configure_buffered_dispatch_profile_counting_service( +rocprofiler_configure_buffered_dispatch_counting_service( rocprofiler_context_id_t context_id, rocprofiler_buffer_id_t buffer_id, - rocprofiler_profile_counting_dispatch_callback_t callback, + rocprofiler_dispatch_counting_service_callback_t callback, void* callback_data_args) { return rocprofiler::counters::configure_buffered_dispatch( @@ -62,9 +62,9 @@ rocprofiler_configure_buffered_dispatch_profile_counting_service( * @return ::rocprofiler_status_t */ rocprofiler_status_t -rocprofiler_configure_callback_dispatch_profile_counting_service( +rocprofiler_configure_callback_dispatch_counting_service( rocprofiler_context_id_t context_id, - rocprofiler_profile_counting_dispatch_callback_t dispatch_callback, + rocprofiler_dispatch_counting_service_callback_t dispatch_callback, void* dispatch_callback_args, rocprofiler_profile_counting_record_callback_t record_callback, void* record_callback_args) diff --git a/source/lib/rocprofiler-sdk/hsa/agent_cache.cpp b/source/lib/rocprofiler-sdk/hsa/agent_cache.cpp index dd4ba1b9a5..2805bc37f9 100644 --- a/source/lib/rocprofiler-sdk/hsa/agent_cache.cpp +++ b/source/lib/rocprofiler-sdk/hsa/agent_cache.cpp @@ -126,7 +126,8 @@ namespace rocprofiler namespace hsa { void -AgentCache::init_agent_profile_queue(const CoreApiTable& api, const AmdExtTable& ext) const +AgentCache::init_device_counting_service_queue(const CoreApiTable& api, + const AmdExtTable& ext) const { static std::mutex m_mutex; std::lock_guard lock(m_mutex); @@ -135,7 +136,7 @@ AgentCache::init_agent_profile_queue(const CoreApiTable& api, const AmdExtTable& const auto* agent_ctx = []() -> const context* { for(auto& ctx : rocprofiler::context::get_registered_contexts()) { - if(ctx->agent_counter_collection) return ctx; + if(ctx->device_counter_collection) return ctx; } return nullptr; }(); @@ -178,7 +179,7 @@ AgentCache::AgentCache(const rocprofiler_agent_t* rocp_agent, { init_cpu_pool(ext_table, *this); init_gpu_pool(ext_table, *this); - init_agent_profile_queue(api, ext_table); + init_device_counting_service_queue(api, ext_table); } catch(std::runtime_error& e) { ROCP_WARNING << fmt::format( diff --git a/source/lib/rocprofiler-sdk/hsa/agent_cache.hpp b/source/lib/rocprofiler-sdk/hsa/agent_cache.hpp index cced1c8abc..7326ce6338 100644 --- a/source/lib/rocprofiler-sdk/hsa/agent_cache.hpp +++ b/source/lib/rocprofiler-sdk/hsa/agent_cache.hpp @@ -73,7 +73,7 @@ public: std::string_view name() const { return m_name; } size_t index() const { return m_index; } - void init_agent_profile_queue(const CoreApiTable& api, const AmdExtTable& ext) const; + void init_device_counting_service_queue(const CoreApiTable& api, const AmdExtTable& ext) const; bool operator==(const rocprofiler_agent_t*) const; bool operator==(hsa_agent_t) const; diff --git a/source/lib/rocprofiler-sdk/hsa/aql_packet.cpp b/source/lib/rocprofiler-sdk/hsa/aql_packet.cpp index 72fec5c53b..79d0b21448 100644 --- a/source/lib/rocprofiler-sdk/hsa/aql_packet.cpp +++ b/source/lib/rocprofiler-sdk/hsa/aql_packet.cpp @@ -21,6 +21,7 @@ // THE SOFTWARE. #include "lib/rocprofiler-sdk/hsa/aql_packet.hpp" +#include #include #include #include "lib/common/logging.hpp" @@ -115,7 +116,19 @@ CounterAQLPacket::CounterAQLPacket(aqlprofile_agent_handle_t ag &CounterMemoryPool::Free, &CounterMemoryPool::Copy, reinterpret_cast(&pool)); - if(status != HSA_STATUS_SUCCESS) ROCP_FATAL << "Could not create PMC packets!"; + if(status != HSA_STATUS_SUCCESS) + { + std::string event_list; + for(const auto& event : events) + { + event_list += fmt::format("[{},{},{}],", + event.block_index, + event.event_id, + static_cast(event.block_name)); + } + ROCP_FATAL << "Could not create PMC packets! AQLProfile Return Code: " << status + << " Events: " << event_list; + } packets.start_packet.header = VENDOR_BIT; packets.stop_packet.header = VENDOR_BIT | BARRIER_BIT; diff --git a/source/lib/rocprofiler-sdk/hsa/queue_controller.cpp b/source/lib/rocprofiler-sdk/hsa/queue_controller.cpp index 57501f5a7f..9f2e33d92c 100644 --- a/source/lib/rocprofiler-sdk/hsa/queue_controller.cpp +++ b/source/lib/rocprofiler-sdk/hsa/queue_controller.cpp @@ -277,7 +277,7 @@ QueueController::init(CoreApiTable& core_table, AmdExtTable& ext_table) itr->is_tracing(ROCPROFILER_BUFFER_TRACING_SCRATCH_MEMORY); if(itr->counter_collection || itr->pc_sampler || has_kernel_tracing || - has_scratch_reporting || itr->agent_counter_collection || itr->agent_thread_trace || + has_scratch_reporting || itr->device_counter_collection || itr->agent_thread_trace || itr->dispatch_thread_trace) { enable_intercepter = true; diff --git a/source/lib/rocprofiler-sdk/pc_sampling/service.cpp b/source/lib/rocprofiler-sdk/pc_sampling/service.cpp index 8e16fde64d..e56c5fe87a 100644 --- a/source/lib/rocprofiler-sdk/pc_sampling/service.cpp +++ b/source/lib/rocprofiler-sdk/pc_sampling/service.cpp @@ -171,7 +171,7 @@ configure_pc_sampling_service(context::context* ctx, // counter collection service can enable clock gating and hang might appear. // As a workaround, PC sampling and (dispatch) counter collection service // cannot coexist in the same context. - if(ctx->counter_collection || ctx->agent_counter_collection) + if(ctx->counter_collection || ctx->device_counter_collection) { return ROCPROFILER_STATUS_ERROR_CONTEXT_CONFLICT; } diff --git a/source/lib/rocprofiler-sdk/pc_sampling/tests/pc_sampling_vs_counter_collection.cpp b/source/lib/rocprofiler-sdk/pc_sampling/tests/pc_sampling_vs_counter_collection.cpp index 79b8610de3..feb9942ebd 100644 --- a/source/lib/rocprofiler-sdk/pc_sampling/tests/pc_sampling_vs_counter_collection.cpp +++ b/source/lib/rocprofiler-sdk/pc_sampling/tests/pc_sampling_vs_counter_collection.cpp @@ -198,7 +198,7 @@ rocprofiler_pc_sampling_callback(rocprofiler_context_id_t /*context_id*/, // =================== Functions related to the counter collection service void -record_callback(rocprofiler_profile_counting_dispatch_data_t /*dispatch_data*/, +record_callback(rocprofiler_dispatch_counting_service_data_t /*dispatch_data*/, rocprofiler_record_counter_t* /*record_data*/, size_t /*record_count*/, rocprofiler_user_data_t /*user_data*/, @@ -206,7 +206,7 @@ record_callback(rocprofiler_profile_counting_dispatch_data_t /*dispatch_data*/, {} void -dispatch_callback(rocprofiler_profile_counting_dispatch_data_t /*dispatch_data*/, +dispatch_callback(rocprofiler_dispatch_counting_service_data_t /*dispatch_data*/, rocprofiler_profile_config_id_t* /*config*/, rocprofiler_user_data_t* /*user_data*/, void* /*callback_data_args*/) @@ -446,7 +446,7 @@ TEST(pc_sampling, pc_sampling_vs_dispatch_counter_collection) auto dispatch_counter_collection_setup_fn = [](rocprofiler_context_id_t context_id, rocprofiler_agent_id_t /*agent_id*/) { // Configure dispatch counter collection service on all agents - EXPECT_EQ(rocprofiler_configure_callback_dispatch_profile_counting_service( + EXPECT_EQ(rocprofiler_configure_callback_dispatch_counting_service( context_id, dispatch_callback, nullptr, record_callback, nullptr), ROCPROFILER_STATUS_SUCCESS); }; @@ -454,10 +454,10 @@ TEST(pc_sampling, pc_sampling_vs_dispatch_counter_collection) pc_sampling_vs_counter_collection(dispatch_counter_collection_setup_fn); } -TEST(pc_sampling, pc_sampling_vs_agent_counter_collection) +TEST(pc_sampling, pc_sampling_vs_device_counter_collection) { - auto agent_counter_collection_setup_fn = [](rocprofiler_context_id_t context_id, - rocprofiler_agent_id_t agent_id) { + auto device_counter_collection_setup_fn = [](rocprofiler_context_id_t context_id, + rocprofiler_agent_id_t agent_id) { rocprofiler_buffer_id_t cc_buf_id; // Create PC sampling buffer EXPECT_EQ(rocprofiler_create_buffer(context_id, @@ -469,12 +469,12 @@ TEST(pc_sampling, pc_sampling_vs_agent_counter_collection) &cc_buf_id), ROCPROFILER_STATUS_SUCCESS); - EXPECT_EQ(rocprofiler_configure_agent_profile_counting_service( + EXPECT_EQ(rocprofiler_configure_device_counting_service( context_id, cc_buf_id, agent_id, set_profile, nullptr), ROCPROFILER_STATUS_SUCCESS); }; - pc_sampling_vs_counter_collection(agent_counter_collection_setup_fn); + pc_sampling_vs_counter_collection(device_counter_collection_setup_fn); } TEST(pc_sampling, dispatch_counter_collection_vs_pc_sampling) @@ -482,7 +482,7 @@ TEST(pc_sampling, dispatch_counter_collection_vs_pc_sampling) auto dispatch_counter_collection_setup_fn = [](rocprofiler_context_id_t context_id, rocprofiler_agent_id_t /*agent_id*/) { // Configure dispatch counter collection service on all agents - EXPECT_EQ(rocprofiler_configure_callback_dispatch_profile_counting_service( + EXPECT_EQ(rocprofiler_configure_callback_dispatch_counting_service( context_id, dispatch_callback, nullptr, record_callback, nullptr), ROCPROFILER_STATUS_ERROR_CONTEXT_CONFLICT); }; @@ -490,10 +490,10 @@ TEST(pc_sampling, dispatch_counter_collection_vs_pc_sampling) counter_collection_vs_pc_sampling(dispatch_counter_collection_setup_fn); } -TEST(pc_sampling, agent_counter_collection_vs_pc_sampling) +TEST(pc_sampling, device_counter_collection_vs_pc_sampling) { - auto agent_counter_collection_setup_fn = [](rocprofiler_context_id_t context_id, - rocprofiler_agent_id_t agent_id) { + auto device_counter_collection_setup_fn = [](rocprofiler_context_id_t context_id, + rocprofiler_agent_id_t agent_id) { rocprofiler_buffer_id_t cc_buf_id; // Create PC sampling buffer EXPECT_EQ(rocprofiler_create_buffer(context_id, @@ -505,10 +505,10 @@ TEST(pc_sampling, agent_counter_collection_vs_pc_sampling) &cc_buf_id), ROCPROFILER_STATUS_SUCCESS); - EXPECT_EQ(rocprofiler_configure_agent_profile_counting_service( + EXPECT_EQ(rocprofiler_configure_device_counting_service( context_id, cc_buf_id, agent_id, set_profile, nullptr), ROCPROFILER_STATUS_ERROR_CONTEXT_CONFLICT); }; - counter_collection_vs_pc_sampling(agent_counter_collection_setup_fn); + counter_collection_vs_pc_sampling(device_counter_collection_setup_fn); } diff --git a/source/lib/rocprofiler-sdk/registration.cpp b/source/lib/rocprofiler-sdk/registration.cpp index 6ca4694bf0..efb7fca394 100644 --- a/source/lib/rocprofiler-sdk/registration.cpp +++ b/source/lib/rocprofiler-sdk/registration.cpp @@ -631,6 +631,7 @@ finalize() std::call_once(_once, []() { set_fini_status(-1); hsa::async_copy_fini(); + counters::device_counting_service_finalize(); hsa::queue_controller_fini(); thread_trace::finalize(); page_migration::finalize(); @@ -786,7 +787,7 @@ rocprofiler_set_api_table(const char* name, rocprofiler::agent::construct_agent_cache(hsa_api_table); rocprofiler::hsa::queue_controller_init(hsa_api_table); // Process agent ctx's that were started prior to HSA init - rocprofiler::counters::agent_profile_hsa_registration(); + rocprofiler::counters::device_counting_service_hsa_registration(); rocprofiler::hsa::async_copy_init(hsa_api_table, lib_instance); rocprofiler::code_object::initialize(hsa_api_table); diff --git a/source/lib/rocprofiler-sdk/tests/details/agent.cpp b/source/lib/rocprofiler-sdk/tests/details/agent.cpp index f923db2dc0..44c3ba1030 100644 --- a/source/lib/rocprofiler-sdk/tests/details/agent.cpp +++ b/source/lib/rocprofiler-sdk/tests/details/agent.cpp @@ -122,7 +122,7 @@ AcquireAgentInfoEntry(hsa_agent_t agent, agent_info_t* agent_i) RET_IF_HSA_ERR(err); // Get profile supported by the agent - err = hsa_agent_get_info(agent, HSA_AGENT_INFO_PROFILE, &agent_i->agent_profile); + err = hsa_agent_get_info(agent, HSA_AGENT_INFO_PROFILE, &agent_i->device_counting_service); RET_IF_HSA_ERR(err); // Get floating-point rounding mode diff --git a/source/lib/rocprofiler-sdk/tests/details/agent.hpp b/source/lib/rocprofiler-sdk/tests/details/agent.hpp index 25644071b7..18bd3bfb7b 100644 --- a/source/lib/rocprofiler-sdk/tests/details/agent.hpp +++ b/source/lib/rocprofiler-sdk/tests/details/agent.hpp @@ -56,28 +56,28 @@ struct system_info_t // calls, and is later used for reference when displaying the information. struct agent_info_t { - hsa_agent_t hsa_agent = {.handle = 0}; - char name[64] = {'\0'}; - char vendor_name[64] = {'\0'}; - char device_mkt_name[64] = {'\0'}; - hsa_agent_feature_t agent_feature = {}; - hsa_profile_t agent_profile = {}; - hsa_default_float_rounding_mode_t float_rounding_mode = {}; - uint32_t max_queue = 0; - uint32_t queue_min_size = 0; - uint32_t queue_max_size = 0; - hsa_queue_type_t queue_type = {}; - uint32_t node = 0; - hsa_device_type_t device_type = {}; - uint32_t cache_size[4] = {0, 0, 0, 0}; - uint32_t chip_id = 0; - uint32_t cacheline_size = 0; - uint32_t max_clock_freq = 0; - uint32_t internal_node_id = 0; - uint32_t max_addr_watch_pts = 0; - uint32_t family_id = 0; - uint32_t ucode_version = 0; - uint32_t sdma_ucode_version = 0; + hsa_agent_t hsa_agent = {.handle = 0}; + char name[64] = {'\0'}; + char vendor_name[64] = {'\0'}; + char device_mkt_name[64] = {'\0'}; + hsa_agent_feature_t agent_feature = {}; + hsa_profile_t device_counting_service = {}; + hsa_default_float_rounding_mode_t float_rounding_mode = {}; + uint32_t max_queue = 0; + uint32_t queue_min_size = 0; + uint32_t queue_max_size = 0; + hsa_queue_type_t queue_type = {}; + uint32_t node = 0; + hsa_device_type_t device_type = {}; + uint32_t cache_size[4] = {0, 0, 0, 0}; + uint32_t chip_id = 0; + uint32_t cacheline_size = 0; + uint32_t max_clock_freq = 0; + uint32_t internal_node_id = 0; + uint32_t max_addr_watch_pts = 0; + uint32_t family_id = 0; + uint32_t ucode_version = 0; + uint32_t sdma_ucode_version = 0; // HSA_AMD_AGENT_INFO_MEMORY_WIDTH is deprecated, so exclude // uint32_t mem_max_freq; Not supported by get_info uint32_t compute_unit = 0; diff --git a/tests/tools/json-tool.cpp b/tests/tools/json-tool.cpp index 781d2bbcea..05a7a0dd12 100644 --- a/tests/tools/json-tool.cpp +++ b/tests/tools/json-tool.cpp @@ -43,7 +43,7 @@ #include #include #include -#include +#include #include #include #include @@ -428,11 +428,11 @@ struct scratch_memory_callback_record_t struct profile_counting_record { - profile_counting_record(rocprofiler_profile_counting_dispatch_record_t hdr) + profile_counting_record(rocprofiler_dispatch_counting_service_record_t hdr) : header{hdr} {} - rocprofiler_profile_counting_dispatch_record_t header = {}; + rocprofiler_dispatch_counting_service_record_t header = {}; std::vector data = {}; profile_counting_record() = default; @@ -501,7 +501,7 @@ set_external_correlation_id(rocprofiler_thread_id_t t } void -dispatch_callback(rocprofiler_profile_counting_dispatch_data_t dispatch_data, +dispatch_callback(rocprofiler_dispatch_counting_service_data_t dispatch_data, rocprofiler_profile_config_id_t* config, rocprofiler_user_data_t* /*user_data*/, void* /*callback_data_args*/) @@ -584,7 +584,7 @@ dispatch_callback(rocprofiler_profile_counting_dispatch_data_t dispatch_data, } // Create a colleciton profile for the counters - rocprofiler_profile_config_id_t profile; + rocprofiler_profile_config_id_t profile = {.handle = 0}; ROCPROFILER_CALL(rocprofiler_create_profile_config(dispatch_data.dispatch_info.agent_id, collect_counters.data(), collect_counters.size(), @@ -839,7 +839,7 @@ tool_tracing_buffered(rocprofiler_context_id_t /*context*/, header->kind == ROCPROFILER_COUNTER_RECORD_PROFILE_COUNTING_DISPATCH_HEADER) { auto* profiler_record = - static_cast(header->payload); + static_cast(header->payload); counter_collection_bf_records.emplace_back(*profiler_record); } else if(header->category == ROCPROFILER_BUFFER_CATEGORY_COUNTERS && @@ -848,7 +848,7 @@ tool_tracing_buffered(rocprofiler_context_id_t /*context*/, auto* profiler_record = static_cast(header->payload); if(counter_collection_bf_records.empty()) throw std::runtime_error{ - "missing rocprofiler_profile_counting_dispatch_record_t (header)"}; + "missing rocprofiler_dispatch_counting_service_record_t (header)"}; counter_collection_bf_records.back().emplace_back(*profiler_record); } else @@ -1303,7 +1303,7 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data) "buffer tracing service for rccl api configure"); ROCPROFILER_CALL( - rocprofiler_configure_buffered_dispatch_profile_counting_service( + rocprofiler_configure_buffered_dispatch_counting_service( counter_collection_ctx, counter_collection_buffer, dispatch_callback, nullptr), "setup buffered service");