From 4cd121e27b3e02d378b12bb0ef829e19ddfc71db Mon Sep 17 00:00:00 2001 From: "Welton, Benjamin" Date: Sun, 23 Mar 2025 23:37:33 -0700 Subject: [PATCH] [SDK] Release 1.0 Public API Modifications (#277) * Make sure all structs/enums can be forward declared * Updates to counter collection - consistency updates and cleanup * Conversion of dimension information to info struct * Added deprecated folder * Testing changes * merge changes * Fix shadowed variable * Source code formatting * Fix shadowed variable * Update rocprofiler_counter_info_v1_t member names * Split version.h into version.h and ext_version.h - ext_version.h contains external version info, e.g. ROCPROFILER_HSA_API_TABLE_MAJOR_VERSION, ROCPROFILER_HSA_RUNTIME_VERSION - this reduces amount of recompilation after a commit since version.h gets updated with the git revision * profile_config -> counter_config * EOF new line * [Samples] Reduce header includes + reorg counter collection samples * Misc compilation fixes - shadowed variables - use of [[deprecated("...")]] in C code - unused variables * Minor misc modifications - use common:: instead of rocprofiler::common:: when inside rocprofiler namespace - counters.cpp - move local anon namespace functions into rocprofiler::counters:: anon namespace - use std::string_view for get_static_string - const ref for get_static_ptr - misc namespace shortening * [Public API] rocprofiler_get_version_triplet + rocprofiler_version_triplet_t - struct rocprofiler_version_triplet_t containing fields for the major, minor, and patch version - public API function: rocprofiler_get_version_triplet - define C++ operators for rocprofiler_version_triplet_t - C++ function compute_version_triplet * [Tests] Improve async-copy-testing test - relax constraints - improve logging * Update counter_config.h doxygen docs * ROCPROFILER_SDK_BETA_COMPAT - ppdef which helps with renaming when set to 1 * Remove spurious include * Fix includes for cxx/version.hpp * Doxygen fixes for rocprofiler_get_version and rocprofiler_get_version_triplet * Public API Experimental Designation - ROCPROFILER_SDK_EXPERIMENTAL added to experimental function - "(experimental)" added to doxygen @brief entries * Fix use of assert instead of static_assert in hip/stream.cpp * Use typedef instead of define for rocprofiler_profile_config_id_t * Use inline rocprofiler_{create,destroy}_profile_config instead of ppdef - added * Doxygen for rocprofiler_{create,destroy}_profile_config * ROCPROFILER_SDK_DEPRECATED_WARNINGS * Temporarily comment out ROCPROFILER_SDK_DEPRECATED_WARNINGS=1 * cmake formatting * Misc variable renaming in samples and tests * Fix declarations of types * Fix hip stream tracing service struct name - rocprofiler_callback_tracing_stream_handle_data_t renamed to rocprofiler_callback_tracing_hip_stream_api_data_t * Rename "HIP_STREAM_API" to "HIP_STREAM" --------- Co-authored-by: Jonathan R. Madsen Co-authored-by: Benjamin Welton --- cmake/rocprofiler_build_settings.cmake | 13 + cmake/rocprofiler_interfaces.cmake | 3 + cmake/rocprofiler_options.cmake | 5 + cmake/rocprofiler_utilities.cmake | 2 +- samples/api_buffered_tracing/client.cpp | 6 - samples/api_callback_tracing/client.cpp | 24 +- samples/code_object_isa_decode/client.cpp | 3 - samples/code_object_tracing/client.cpp | 3 - samples/counter_collection/CMakeLists.txt | 14 +- .../{client.cpp => buffered_client.cpp} | 56 ++--- .../counter_collection/callback_client.cpp | 24 +- ...g.cpp => device_counting_async_client.cpp} | 33 ++- ...us.cpp => device_counting_sync_client.cpp} | 78 ++---- ...ization.cpp => device_serialized_main.cpp} | 0 ...p => print_functional_counters_client.cpp} | 54 ++--- .../client.cpp | 8 - samples/openmp_target/client.cpp | 5 - samples/pc_sampling/client.cpp | 3 - .../counter_collection_services.rst | 24 +- source/docs/rocprofiler-sdk.dox.in | 15 +- source/include/rocprofiler-sdk/CMakeLists.txt | 13 +- .../amd_detail/thread_trace_core.h | 4 +- .../amd_detail/thread_trace_dispatch.h | 2 +- .../include/rocprofiler-sdk/buffer_tracing.h | 18 +- .../rocprofiler-sdk/callback_tracing.h | 24 +- .../{profile_config.h => counter_config.h} | 38 +-- source/include/rocprofiler-sdk/counters.h | 132 +++++----- .../rocprofiler-sdk/cxx/CMakeLists.txt | 2 +- source/include/rocprofiler-sdk/cxx/hash.hpp | 2 +- .../include/rocprofiler-sdk/cxx/operators.hpp | 43 +++- .../rocprofiler-sdk/cxx/serialization.hpp | 24 ++ .../include/rocprofiler-sdk/cxx/version.hpp | 59 +++++ source/include/rocprofiler-sdk/defines.h | 78 ++++++ .../rocprofiler-sdk/deprecated/CMakeLists.txt | 9 + .../rocprofiler-sdk/deprecated/counters.h | 87 +++++++ .../deprecated/profile_config.h | 56 +++++ .../rocprofiler-sdk/device_counting_service.h | 66 ++--- .../dispatch_counting_service.h | 62 ++--- .../rocprofiler-sdk/experimental/counters.h | 9 +- .../include/rocprofiler-sdk/ext_version.h.in | 55 +++++ .../rocprofiler-sdk/external_correlation.h | 29 ++- source/include/rocprofiler-sdk/fwd.h | 131 +++++----- source/include/rocprofiler-sdk/hip/api_args.h | 1 - .../rocprofiler-sdk/hip/compiler_api_id.h | 4 +- .../rocprofiler-sdk/hip/runtime_api_id.h | 4 +- source/include/rocprofiler-sdk/hip/table_id.h | 2 +- .../rocprofiler-sdk/hsa/amd_ext_api_id.h | 4 +- source/include/rocprofiler-sdk/hsa/api_args.h | 2 +- .../rocprofiler-sdk/hsa/api_trace_version.h | 2 +- .../include/rocprofiler-sdk/hsa/core_api_id.h | 4 +- .../rocprofiler-sdk/hsa/finalize_ext_api_id.h | 4 +- .../rocprofiler-sdk/hsa/image_ext_api_id.h | 4 +- .../rocprofiler-sdk/hsa/scratch_memory_args.h | 6 +- .../rocprofiler-sdk/hsa/scratch_memory_id.h | 2 +- source/include/rocprofiler-sdk/hsa/table_id.h | 2 +- .../include/rocprofiler-sdk/intercept_table.h | 40 ++-- .../rocprofiler-sdk/internal_threading.h | 42 ++-- .../rocprofiler-sdk/kfd/page_migration_args.h | 2 +- .../rocprofiler-sdk/kfd/page_migration_id.h | 7 +- .../include/rocprofiler-sdk/marker/api_args.h | 1 - .../include/rocprofiler-sdk/marker/api_id.h | 6 +- .../include/rocprofiler-sdk/marker/table_id.h | 2 +- source/include/rocprofiler-sdk/ompt.h | 28 +++ .../include/rocprofiler-sdk/ompt/api_args.h | 1 - source/include/rocprofiler-sdk/ompt/api_id.h | 2 +- source/include/rocprofiler-sdk/pc_sampling.h | 91 +++---- .../include/rocprofiler-sdk/rccl/api_args.h | 1 - source/include/rocprofiler-sdk/rccl/api_id.h | 6 +- .../include/rocprofiler-sdk/rccl/table_id.h | 2 +- source/include/rocprofiler-sdk/registration.h | 29 +-- .../rocprofiler-sdk/rocdecode/CMakeLists.txt | 2 +- .../rocprofiler-sdk/rocdecode/api_args.h | 2 +- .../rocprofiler-sdk/rocdecode/api_id.h | 4 +- .../rocprofiler-sdk/rocdecode/api_trace.h | 31 +++ .../rocprofiler-sdk/rocdecode/table_id.h | 2 +- .../rocprofiler-sdk/rocjpeg/api_args.h | 1 - .../include/rocprofiler-sdk/rocjpeg/api_id.h | 4 +- .../rocprofiler-sdk/rocjpeg/table_id.h | 3 +- source/include/rocprofiler-sdk/rocprofiler.h | 28 ++- source/include/rocprofiler-sdk/spm.h | 5 +- source/include/rocprofiler-sdk/version.h.in | 25 +- source/lib/common/abi.hpp | 2 +- source/lib/output/counter_info.hpp | 6 +- source/lib/output/metadata.cpp | 37 +-- source/lib/rocprofiler-sdk-roctx/abi.cpp | 2 +- source/lib/rocprofiler-sdk-roctx/roctx.cpp | 2 +- source/lib/rocprofiler-sdk-tool/tool.cpp | 39 ++- source/lib/rocprofiler-sdk/CMakeLists.txt | 8 +- source/lib/rocprofiler-sdk/buffer_tracing.cpp | 6 +- .../lib/rocprofiler-sdk/callback_tracing.cpp | 8 +- ...{profile_config.cpp => counter_config.cpp} | 12 +- source/lib/rocprofiler-sdk/counters.cpp | 226 ++++++++++++++---- .../rocprofiler-sdk/counters/controller.cpp | 41 ++-- .../rocprofiler-sdk/counters/controller.hpp | 42 ++-- source/lib/rocprofiler-sdk/counters/core.cpp | 34 +-- source/lib/rocprofiler-sdk/counters/core.hpp | 44 ++-- .../counters/device_counting.cpp | 8 +- .../counters/device_counting.hpp | 6 +- .../counters/dispatch_handlers.cpp | 4 +- .../counters/sample_processing.hpp | 2 +- .../rocprofiler-sdk/counters/tests/core.cpp | 36 +-- .../counters/tests/device_counting.cpp | 10 +- .../counters/tests/dimension.cpp | 5 + .../counters/tests/metrics_test.cpp | 66 ++++- .../device_counting_service.cpp | 10 +- ...file.cpp => dispatch_counting_service.cpp} | 20 +- source/lib/rocprofiler-sdk/hip/abi.cpp | 2 +- .../rocprofiler-sdk/hip/details/format.hpp | 1 - source/lib/rocprofiler-sdk/hip/stream.cpp | 38 +-- source/lib/rocprofiler-sdk/hip/utils.hpp | 2 +- source/lib/rocprofiler-sdk/hsa/abi.cpp | 2 +- source/lib/rocprofiler-sdk/hsa/queue.hpp | 20 +- .../rocprofiler-sdk/hsa/scratch_memory.cpp | 3 + source/lib/rocprofiler-sdk/hsa/utils.hpp | 2 +- source/lib/rocprofiler-sdk/marker/utils.hpp | 2 +- source/lib/rocprofiler-sdk/ompt/utils.hpp | 1 - .../pc_sampling_vs_counter_collection.cpp | 4 +- source/lib/rocprofiler-sdk/rccl/abi.cpp | 2 +- source/lib/rocprofiler-sdk/registration.cpp | 16 +- source/lib/rocprofiler-sdk/rocdecode/abi.cpp | 2 +- source/lib/rocprofiler-sdk/rocjpeg/abi.cpp | 2 +- source/lib/rocprofiler-sdk/rocprofiler.cpp | 9 + .../lib/rocprofiler-sdk/tests/hsa_barrier.cpp | 10 +- source/lib/rocprofiler-sdk/tests/version.cpp | 3 + .../rocprofiler-avail/rocprofv3_avail.cpp | 15 +- tests/async-copy-tracing/validate.py | 71 +++++- tests/tools/json-tool.cpp | 10 +- 127 files changed, 1629 insertions(+), 988 deletions(-) rename samples/counter_collection/{client.cpp => buffered_client.cpp} (92%) rename samples/counter_collection/{device_counting.cpp => device_counting_async_client.cpp} (95%) rename samples/counter_collection/{device_counting_synchronous.cpp => device_counting_sync_client.cpp} (88%) rename samples/counter_collection/{per_dev_serialization.cpp => device_serialized_main.cpp} (100%) rename samples/counter_collection/{print_functional_counters.cpp => print_functional_counters_client.cpp} (88%) rename source/include/rocprofiler-sdk/{profile_config.h => counter_config.h} (66%) create mode 100644 source/include/rocprofiler-sdk/cxx/version.hpp create mode 100644 source/include/rocprofiler-sdk/deprecated/CMakeLists.txt create mode 100644 source/include/rocprofiler-sdk/deprecated/counters.h create mode 100644 source/include/rocprofiler-sdk/deprecated/profile_config.h create mode 100644 source/include/rocprofiler-sdk/ext_version.h.in create mode 100644 source/include/rocprofiler-sdk/rocdecode/api_trace.h rename source/lib/rocprofiler-sdk/{profile_config.cpp => counter_config.cpp} (91%) rename source/lib/rocprofiler-sdk/{dispatch_profile.cpp => dispatch_counting_service.cpp} (81%) diff --git a/cmake/rocprofiler_build_settings.cmake b/cmake/rocprofiler_build_settings.cmake index f1ddc2e7e0..3abdbed653 100644 --- a/cmake/rocprofiler_build_settings.cmake +++ b/cmake/rocprofiler_build_settings.cmake @@ -203,6 +203,19 @@ if(ROCPROFILER_BUILD_CI_STRICT_TIMESTAMPS) INTERFACE ROCPROFILER_CI_STRICT_TIMESTAMPS) endif() +# ----------------------------------------------------------------------------------------# +# extra flags for compiling with experimental warnings +# +target_compile_definitions(rocprofiler-sdk-experimental-flags + INTERFACE ROCPROFILER_SDK_EXPERIMENTAL_WARNINGS=1) +rocprofiler_target_compile_options(rocprofiler-sdk-experimental-flags + INTERFACE "-Wno-deprecated-declarations") + +if(ROCPROFILER_BUILD_EXPERIMENTAL_WARNINGS) + target_link_libraries(rocprofiler-sdk-build-flags + INTERFACE rocprofiler-sdk::rocprofiler-sdk-experimental-flags) +endif() + # ----------------------------------------------------------------------------------------# # user customization # diff --git a/cmake/rocprofiler_interfaces.cmake b/cmake/rocprofiler_interfaces.cmake index 8f84d6e098..45b15042c8 100644 --- a/cmake/rocprofiler_interfaces.cmake +++ b/cmake/rocprofiler_interfaces.cmake @@ -38,6 +38,9 @@ rocprofiler_add_interface_library(rocprofiler-sdk-release-flags rocprofiler_add_interface_library(rocprofiler-sdk-stack-protector "Adds stack-protector compiler flags" INTERNAL) rocprofiler_add_interface_library(rocprofiler-sdk-memcheck INTERFACE INTERNAL) +rocprofiler_add_interface_library( + rocprofiler-sdk-experimental-flags + "Compiler flags for experimental feature compilation" INTERNAL) # # interfaces for libraries (general) diff --git a/cmake/rocprofiler_options.cmake b/cmake/rocprofiler_options.cmake index 05b960694b..2da7445616 100644 --- a/cmake/rocprofiler_options.cmake +++ b/cmake/rocprofiler_options.cmake @@ -81,6 +81,11 @@ rocprofiler_add_option(ROCPROFILER_UNSAFE_NO_VERSION_CHECK rocprofiler_add_option( ROCPROFILER_REGENERATE_COUNTERS_PARSER "Regenerate the counter parser (requires bison and flex)" OFF ADVANCED) +rocprofiler_add_option( + ROCPROFILER_BUILD_EXPERIMENTAL_WARNINGS + "Enable warnings for experimental features but hide with -Wno-deprecated-declarations (this ensures that experimental warning message does not break macros)" + OFF + ADVANCED) # In the future, we will do this even with clang-tidy enabled foreach(_OPT ROCPROFILER_BUILD_DEVELOPER ROCPROFILER_BUILD_WERROR) diff --git a/cmake/rocprofiler_utilities.cmake b/cmake/rocprofiler_utilities.cmake index 344c5b61c4..4bd5e7062b 100644 --- a/cmake/rocprofiler_utilities.cmake +++ b/cmake/rocprofiler_utilities.cmake @@ -968,7 +968,7 @@ function(rocprofiler_parse_hsa_api_table_versions _TARGET) _VAL "${_LINE}") - # used with cmakedefine in source/include/rocprofiler-sdk/version.h.in + # used with cmakedefine in source/include/rocprofiler-sdk/ext_version.h.in if(_VAR AND _VAL) set(ROCPROFILER_${_VAR} "${_VAL}" diff --git a/samples/api_buffered_tracing/client.cpp b/samples/api_buffered_tracing/client.cpp index 81ed12cf5d..cfe38f6149 100644 --- a/samples/api_buffered_tracing/client.cpp +++ b/samples/api_buffered_tracing/client.cpp @@ -33,12 +33,6 @@ #include "client.hpp" -#include -#include -#include -#include -#include -#include #include #include diff --git a/samples/api_callback_tracing/client.cpp b/samples/api_callback_tracing/client.cpp index 57d6e1e537..af866bdb02 100644 --- a/samples/api_callback_tracing/client.cpp +++ b/samples/api_callback_tracing/client.cpp @@ -33,11 +33,10 @@ #include "client.hpp" -#include -#include -#include #include #include +#include +#include #include "common/call_stack.hpp" #include "common/defines.hpp" @@ -318,7 +317,7 @@ stop() extern "C" rocprofiler_tool_configure_result_t* rocprofiler_configure(uint32_t version, - const char* runtime_version, + const char* version_string, uint32_t priority, rocprofiler_client_id_t* id) { @@ -336,21 +335,26 @@ rocprofiler_configure(uint32_t version, // generate info string auto info = std::stringstream{}; info << id->name << " (priority=" << priority << ") is using rocprofiler-sdk v" << major << "." - << minor << "." << patch << " (" << runtime_version << ")"; + << minor << "." << patch << " (" << version_string << ")"; std::clog << info.str() << std::endl; // demonstration of alternative way to get the version info { - auto version_info = std::array{}; - ROCPROFILER_CALL( - rocprofiler_get_version(&version_info.at(0), &version_info.at(1), &version_info.at(2)), - "failed to get version info"); + auto runtime_version = rocprofiler_version_triplet_t{}; + ROCPROFILER_CALL(rocprofiler_get_version_triplet(&runtime_version), + "failed to get version info"); - if(std::array{major, minor, patch} != version_info) + if(rocprofiler_version_triplet_t{major, minor, patch} != runtime_version) { throw std::runtime_error{"version info mismatch"}; } + + if(rocprofiler_version_triplet_t{major, minor, patch} != + rocprofiler::sdk::version::compute_version_triplet<100>(version)) + { + throw std::runtime_error{"version triplet incorrectly calculated"}; + } } // data passed around all the callbacks diff --git a/samples/code_object_isa_decode/client.cpp b/samples/code_object_isa_decode/client.cpp index 630b78fe5a..12264a0927 100644 --- a/samples/code_object_isa_decode/client.cpp +++ b/samples/code_object_isa_decode/client.cpp @@ -33,9 +33,6 @@ * @brief Example rocprofiler client (tool) */ -#include -#include -#include #include #include #include diff --git a/samples/code_object_tracing/client.cpp b/samples/code_object_tracing/client.cpp index 98c2394a06..7186ad1a18 100644 --- a/samples/code_object_tracing/client.cpp +++ b/samples/code_object_tracing/client.cpp @@ -31,9 +31,6 @@ * @brief Example rocprofiler client (tool) */ -#include -#include -#include #include #include diff --git a/samples/counter_collection/CMakeLists.txt b/samples/counter_collection/CMakeLists.txt index ecf35bbbf3..3f124e0bd2 100644 --- a/samples/counter_collection/CMakeLists.txt +++ b/samples/counter_collection/CMakeLists.txt @@ -34,7 +34,7 @@ endif() find_package(rocprofiler-sdk REQUIRED) add_library(counter-collection-buffer-client SHARED) -target_sources(counter-collection-buffer-client PRIVATE client.cpp client.hpp) +target_sources(counter-collection-buffer-client PRIVATE buffered_client.cpp client.hpp) target_link_libraries( counter-collection-buffer-client PUBLIC rocprofiler-sdk::samples-build-flags @@ -59,10 +59,10 @@ set_tests_properties( "${counter-collection-buffer-env}" FAIL_REGULAR_EXPRESSION "${ROCPROFILER_DEFAULT_FAIL_REGEX}") -set_source_files_properties(per_dev_serialization.cpp PROPERTIES LANGUAGE HIP) +set_source_files_properties(device_serialized_main.cpp PROPERTIES LANGUAGE HIP) add_executable(counter-collection-buffer-device-serialization) target_sources(counter-collection-buffer-device-serialization - PRIVATE per_dev_serialization.cpp) + PRIVATE device_serialized_main.cpp) target_link_libraries(counter-collection-buffer-device-serialization PRIVATE counter-collection-buffer-client Threads::Threads) @@ -109,7 +109,7 @@ set_tests_properties( add_library(counter-collection-functional-counter-client SHARED) target_sources(counter-collection-functional-counter-client - PRIVATE print_functional_counters.cpp client.hpp) + PRIVATE print_functional_counters_client.cpp client.hpp) target_link_libraries( counter-collection-functional-counter-client PUBLIC rocprofiler-sdk::samples-build-flags @@ -136,8 +136,8 @@ set_tests_properties( "${ROCPROFILER_DEFAULT_FAIL_REGEX}") add_library(counter-collection-device-profiling-client SHARED) -target_sources(counter-collection-device-profiling-client PRIVATE device_counting.cpp - client.hpp) +target_sources(counter-collection-device-profiling-client + PRIVATE device_counting_async_client.cpp client.hpp) target_link_libraries( counter-collection-device-profiling-client PUBLIC rocprofiler-sdk::samples-build-flags @@ -171,7 +171,7 @@ set_tests_properties( add_library(counter-collection-device-profiling-sync-client SHARED) target_sources(counter-collection-device-profiling-sync-client - PRIVATE device_counting_synchronous.cpp client.hpp) + PRIVATE device_counting_sync_client.cpp client.hpp) target_link_libraries( counter-collection-device-profiling-sync-client PUBLIC rocprofiler-sdk::samples-build-flags diff --git a/samples/counter_collection/client.cpp b/samples/counter_collection/buffered_client.cpp similarity index 92% rename from samples/counter_collection/client.cpp rename to samples/counter_collection/buffered_client.cpp index 5d8af00edf..ba011943e4 100644 --- a/samples/counter_collection/client.cpp +++ b/samples/counter_collection/buffered_client.cpp @@ -22,6 +22,9 @@ #include "client.hpp" +#include +#include + #include #include #include @@ -33,10 +36,6 @@ #include #include -#include -#include -#include - #define ROCPROFILER_CALL(result, msg) \ { \ rocprofiler_status_t CHECKSTATUS = result; \ @@ -104,22 +103,15 @@ fill_dimension_cache(rocprofiler_counter_id_t counter) { assert(*dimension_cache() != nullptr); std::vector dims; - rocprofiler_available_dimensions_cb_t cb = - [](rocprofiler_counter_id_t, - const rocprofiler_record_dimension_info_t* dim_info, - size_t num_dims, - void* user_data) { - std::vector* vec = - static_cast*>(user_data); - for(size_t i = 0; i < num_dims; i++) - { - vec->push_back(dim_info[i]); - } - return ROCPROFILER_STATUS_SUCCESS; - }; - ROCPROFILER_CALL(rocprofiler_iterate_counter_dimensions(counter, cb, &dims), - "Could not iterate counter dimensions"); - (*dimension_cache())->emplace(counter.handle, dims); + rocprofiler_counter_info_v1_t info; + ROCPROFILER_CALL(rocprofiler_query_counter_info( + counter, ROCPROFILER_COUNTER_INFO_VERSION_1, static_cast(&info)), + "Could not query info for counter"); + + (*dimension_cache()) + ->emplace(counter.handle, + std::vector{ + info.dimensions, info.dimensions + info.dimensions_count}); } /** @@ -187,22 +179,22 @@ buffered_callback(rocprofiler_context_id_t, * to select the profile config (and in turn counters) to use when a kernel dispatch * is received. */ -std::unordered_map& +std::unordered_map& get_profile_cache() { - static std::unordered_map profile_cache; + static std::unordered_map profile_cache; return profile_cache; } /** * Callback from rocprofiler when an kernel dispatch is enqueued into the HSA queue. - * rocprofiler_profile_config_id_t* is a return to specify what counters to collect + * rocprofiler_counter_config_id_t* is a return to specify what counters to collect * for this dispatch (dispatch_packet). This example function creates a profile * to collect the counter SQ_WAVES for all kernel dispatch packets. */ void dispatch_callback(rocprofiler_dispatch_counting_service_data_t dispatch_data, - rocprofiler_profile_config_id_t* config, + rocprofiler_counter_config_id_t* config, rocprofiler_user_data_t* /*user_data*/, void* /*callback_data_args*/) { @@ -236,7 +228,7 @@ dispatch_callback(rocprofiler_dispatch_counting_service_data_t dispatch_data, * to consturct them once in advance (i.e. in tool_init()) since there are non-trivial * costs associated with constructing the profile. */ -rocprofiler_profile_config_id_t +rocprofiler_counter_config_id_t build_profile_for_agent(rocprofiler_agent_id_t agent, const std::set& counters_to_collect) { @@ -264,22 +256,22 @@ build_profile_for_agent(rocprofiler_agent_id_t agent, std::vector collect_counters; for(auto& counter : gpu_counters) { - rocprofiler_counter_info_v0_t version; + rocprofiler_counter_info_v0_t info; ROCPROFILER_CALL( rocprofiler_query_counter_info( - counter, ROCPROFILER_COUNTER_INFO_VERSION_0, static_cast(&version)), + counter, ROCPROFILER_COUNTER_INFO_VERSION_0, static_cast(&info)), "Could not query info for counter"); - if(counters_to_collect.count(std::string(version.name)) > 0) + if(counters_to_collect.count(std::string(info.name)) > 0) { - std::clog << "Counter: " << counter.handle << " " << version.name << "\n"; + std::clog << "Counter: " << counter.handle << " " << info.name << "\n"; collect_counters.push_back(counter); fill_dimension_cache(counter); } } // Create and return the profile - rocprofiler_profile_config_id_t profile = {.handle = 0}; - ROCPROFILER_CALL(rocprofiler_create_profile_config( + rocprofiler_counter_config_id_t profile = {.handle = 0}; + ROCPROFILER_CALL(rocprofiler_create_counter_config( agent, collect_counters.data(), collect_counters.size(), &profile), "Could not construct profile cfg"); @@ -375,7 +367,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_counting_service( + ROCPROFILER_CALL(rocprofiler_configure_buffer_dispatch_counting_service( get_client_ctx(), get_buffer(), dispatch_callback, nullptr), "Could not setup buffered service"); diff --git a/samples/counter_collection/callback_client.cpp b/samples/counter_collection/callback_client.cpp index 6120221ddc..0b5321096b 100644 --- a/samples/counter_collection/callback_client.cpp +++ b/samples/counter_collection/callback_client.cpp @@ -22,6 +22,9 @@ #include "client.hpp" +#include +#include + #include #include #include @@ -32,9 +35,6 @@ #include #include -#include -#include - #define ROCPROFILER_CALL(result, msg) \ { \ rocprofiler_status_t CHECKSTATUS = result; \ @@ -96,13 +96,13 @@ record_callback(rocprofiler_dispatch_counting_service_data_t dispatch_data, /** * Callback from rocprofiler when an kernel dispatch is enqueued into the HSA queue. - * rocprofiler_profile_config_id_t* is a return to specify what counters to collect + * rocprofiler_counter_config_id_t* is a return to specify what counters to collect * for this dispatch (dispatch_packet). This example function creates a profile * to collect the counter SQ_WAVES for all kernel dispatch packets. */ void dispatch_callback(rocprofiler_dispatch_counting_service_data_t dispatch_data, - rocprofiler_profile_config_id_t* config, + rocprofiler_counter_config_id_t* config, rocprofiler_user_data_t* /*user_data*/, void* /*callback_data_args*/) { @@ -114,7 +114,7 @@ dispatch_callback(rocprofiler_dispatch_counting_service_data_t dispatch_data, * set. */ static std::shared_mutex m_mutex = {}; - static std::unordered_map profile_cache = {}; + static std::unordered_map profile_cache = {}; auto search_cache = [&]() { if(auto pos = profile_cache.find(dispatch_data.dispatch_info.agent_id.handle); @@ -161,21 +161,21 @@ dispatch_callback(rocprofiler_dispatch_counting_service_data_t dispatch_data, // Look for the counters contained in counters_to_collect in gpu_counters for(auto& counter : gpu_counters) { - rocprofiler_counter_info_v0_t version; + rocprofiler_counter_info_v0_t info; ROCPROFILER_CALL( rocprofiler_query_counter_info( - counter, ROCPROFILER_COUNTER_INFO_VERSION_0, static_cast(&version)), + counter, ROCPROFILER_COUNTER_INFO_VERSION_0, static_cast(&info)), "Could not query info"); - if(counters_to_collect.count(std::string(version.name)) > 0) + if(counters_to_collect.count(std::string(info.name)) > 0) { - std::clog << "Counter: " << counter.handle << " " << version.name << "\n"; + std::clog << "Counter: " << counter.handle << " " << info.name << "\n"; collect_counters.push_back(counter); } } // Create a colleciton profile for the counters - rocprofiler_profile_config_id_t profile = {.handle = 0}; - ROCPROFILER_CALL(rocprofiler_create_profile_config(dispatch_data.dispatch_info.agent_id, + rocprofiler_counter_config_id_t profile = {.handle = 0}; + ROCPROFILER_CALL(rocprofiler_create_counter_config(dispatch_data.dispatch_info.agent_id, collect_counters.data(), collect_counters.size(), &profile), diff --git a/samples/counter_collection/device_counting.cpp b/samples/counter_collection/device_counting_async_client.cpp similarity index 95% rename from samples/counter_collection/device_counting.cpp rename to samples/counter_collection/device_counting_async_client.cpp index 5d2cd6e06c..0d2077d1f8 100644 --- a/samples/counter_collection/device_counting.cpp +++ b/samples/counter_collection/device_counting_async_client.cpp @@ -22,6 +22,9 @@ #include "client.hpp" +#include +#include + #include #include #include @@ -35,10 +38,6 @@ #include #include -#include -#include -#include - #define ROCPROFILER_CALL(result, msg) \ { \ rocprofiler_status_t CHECKSTATUS = result; \ @@ -128,23 +127,23 @@ buffered_callback(rocprofiler_context_id_t, *output_stream << "[" << __FUNCTION__ << "] " << ss.str() << "\n"; } -std::unordered_map& +std::unordered_map& get_profile_cache() { - static std::unordered_map profile_cache; + static std::unordered_map profile_cache; return profile_cache; } /** * Callback from rocprofiler when an kernel dispatch is enqueued into the HSA queue. - * rocprofiler_profile_config_id_t* is a return to specify what counters to collect + * rocprofiler_counter_config_id_t* is a return to specify what counters to collect * for this dispatch (dispatch_packet). This example function creates a profile * to collect the counter SQ_WAVES for all kernel dispatch packets. */ void -set_profile(rocprofiler_context_id_t context_id, - rocprofiler_agent_id_t agent, - rocprofiler_agent_set_profile_callback_t set_config, +set_profile(rocprofiler_context_id_t context_id, + rocprofiler_agent_id_t agent, + rocprofiler_device_counting_agent_cb_t set_config, void*) { /** @@ -170,7 +169,7 @@ set_profile(rocprofiler_context_id_t context_id, } } -rocprofiler_profile_config_id_t +rocprofiler_counter_config_id_t build_profile_for_agent(rocprofiler_agent_id_t agent) { std::set counters_to_collect = {"SQ_WAVES"}; @@ -196,20 +195,20 @@ build_profile_for_agent(rocprofiler_agent_id_t agent) std::vector collect_counters; for(auto& counter : gpu_counters) { - rocprofiler_counter_info_v0_t version; + rocprofiler_counter_info_v0_t info; ROCPROFILER_CALL( rocprofiler_query_counter_info( - counter, ROCPROFILER_COUNTER_INFO_VERSION_0, static_cast(&version)), + counter, ROCPROFILER_COUNTER_INFO_VERSION_0, static_cast(&info)), "Could not query info for counter"); - if(counters_to_collect.count(std::string(version.name)) > 0) + if(counters_to_collect.count(std::string(info.name)) > 0) { - std::clog << "Counter: " << counter.handle << " " << version.name << "\n"; + std::clog << "Counter: " << counter.handle << " " << info.name << "\n"; collect_counters.push_back(counter); } } - rocprofiler_profile_config_id_t profile = {.handle = 0}; - ROCPROFILER_CALL(rocprofiler_create_profile_config( + rocprofiler_counter_config_id_t profile = {.handle = 0}; + ROCPROFILER_CALL(rocprofiler_create_counter_config( agent, collect_counters.data(), collect_counters.size(), &profile), "Could not construct profile cfg"); diff --git a/samples/counter_collection/device_counting_synchronous.cpp b/samples/counter_collection/device_counting_sync_client.cpp similarity index 88% rename from samples/counter_collection/device_counting_synchronous.cpp rename to samples/counter_collection/device_counting_sync_client.cpp index 59ce8dfabf..ab386d55af 100644 --- a/samples/counter_collection/device_counting_synchronous.cpp +++ b/samples/counter_collection/device_counting_sync_client.cpp @@ -22,6 +22,9 @@ #include "client.hpp" +#include +#include + #include #include #include @@ -38,12 +41,6 @@ #include #include -#include -#include -#include -#include -#include - #define ROCPROFILER_CALL(result, msg) \ { \ rocprofiler_status_t CHECKSTATUS = result; \ @@ -99,15 +96,14 @@ private: rocprofiler_agent_id_t agent_ = {}; rocprofiler_context_id_t ctx_ = {}; rocprofiler_buffer_id_t buf_ = {}; - rocprofiler_profile_config_id_t profile_ = {.handle = 0}; + rocprofiler_counter_config_id_t profile_ = {.handle = 0}; - std::map, rocprofiler_profile_config_id_t> cached_profiles_; + std::map, rocprofiler_counter_config_id_t> cached_profiles_; std::map profile_sizes_; mutable std::map id_to_name_; // Internal function used to set the profile for the agent when start_context is called - void set_profile(rocprofiler_context_id_t ctx, - rocprofiler_agent_set_profile_callback_t cb) const; + void set_profile(rocprofiler_context_id_t ctx, rocprofiler_device_counting_agent_cb_t cb) const; // Get the size of a counter in number of records static size_t get_counter_size(rocprofiler_counter_id_t counter); @@ -153,8 +149,8 @@ counter_sampler::counter_sampler(rocprofiler_agent_id_t agent) agent, [](rocprofiler_context_id_t context_id, rocprofiler_agent_id_t, - rocprofiler_agent_set_profile_callback_t set_config, - void* user_data) { + rocprofiler_device_counting_agent_cb_t set_config, + void* user_data) { if(user_data) { auto* sampler = static_cast(user_data); @@ -212,7 +208,7 @@ counter_sampler::sample_counter_values(const std::vector& if(profile_cached == cached_profiles_.end()) { size_t expected_size = 0; - rocprofiler_profile_config_id_t profile = {}; + rocprofiler_counter_config_id_t profile = {}; std::vector gpu_counters; auto roc_counters = get_supported_counters(agent_); for(const auto& counter : counters) @@ -226,7 +222,7 @@ counter_sampler::sample_counter_values(const std::vector& gpu_counters.push_back(it->second); expected_size += get_counter_size(it->second); } - ROCPROFILER_CALL(rocprofiler_create_profile_config( + ROCPROFILER_CALL(rocprofiler_create_counter_config( agent_, gpu_counters.data(), gpu_counters.size(), &profile), "Could not create profile"); cached_profiles_.emplace(counters, profile); @@ -281,8 +277,8 @@ counter_sampler::get_available_agents() } void -counter_sampler::set_profile(rocprofiler_context_id_t ctx, - rocprofiler_agent_set_profile_callback_t cb) const +counter_sampler::set_profile(rocprofiler_context_id_t ctx, + rocprofiler_device_counting_agent_cb_t cb) const { if(profile_.handle != 0) { @@ -293,22 +289,11 @@ counter_sampler::set_profile(rocprofiler_context_id_t ctx, size_t counter_sampler::get_counter_size(rocprofiler_counter_id_t counter) { - size_t size = 1; - rocprofiler_iterate_counter_dimensions( - counter, - [](rocprofiler_counter_id_t, - const rocprofiler_record_dimension_info_t* dim_info, - size_t num_dims, - void* user_data) { - size_t* s = static_cast(user_data); - for(size_t i = 0; i < num_dims; i++) - { - *s *= dim_info[i].instance_size; - } - return ROCPROFILER_STATUS_SUCCESS; - }, - static_cast(&size)); - return size; + rocprofiler_counter_info_v1_t info; + ROCPROFILER_CALL(rocprofiler_query_counter_info( + counter, ROCPROFILER_COUNTER_INFO_VERSION_1, static_cast(&info)), + "Could not query info for counter"); + return info.instance_ids_count; } std::unordered_map @@ -335,12 +320,12 @@ counter_sampler::get_supported_counters(rocprofiler_agent_id_t agent) "Could not fetch supported counters"); for(auto& counter : gpu_counters) { - rocprofiler_counter_info_v0_t version; + rocprofiler_counter_info_v0_t info; ROCPROFILER_CALL( rocprofiler_query_counter_info( - counter, ROCPROFILER_COUNTER_INFO_VERSION_0, static_cast(&version)), + counter, ROCPROFILER_COUNTER_INFO_VERSION_0, static_cast(&info)), "Could not query info for counter"); - out.emplace(version.name, counter); + out.emplace(info.name, counter); } return out; } @@ -348,23 +333,12 @@ counter_sampler::get_supported_counters(rocprofiler_agent_id_t agent) std::vector counter_sampler::get_counter_dimensions(rocprofiler_counter_id_t counter) { - std::vector dims; - rocprofiler_available_dimensions_cb_t cb = - [](rocprofiler_counter_id_t, - const rocprofiler_record_dimension_info_t* dim_info, - size_t num_dims, - void* user_data) { - std::vector* vec = - static_cast*>(user_data); - for(size_t i = 0; i < num_dims; i++) - { - vec->push_back(dim_info[i]); - } - return ROCPROFILER_STATUS_SUCCESS; - }; - ROCPROFILER_CALL(rocprofiler_iterate_counter_dimensions(counter, cb, &dims), - "Could not iterate counter dimensions"); - return dims; + rocprofiler_counter_info_v1_t info; + ROCPROFILER_CALL(rocprofiler_query_counter_info( + counter, ROCPROFILER_COUNTER_INFO_VERSION_1, static_cast(&info)), + "Could not query info for counter"); + return std::vector{ + info.dimensions, info.dimensions + info.dimensions_count}; } std::atomic& diff --git a/samples/counter_collection/per_dev_serialization.cpp b/samples/counter_collection/device_serialized_main.cpp similarity index 100% rename from samples/counter_collection/per_dev_serialization.cpp rename to samples/counter_collection/device_serialized_main.cpp diff --git a/samples/counter_collection/print_functional_counters.cpp b/samples/counter_collection/print_functional_counters_client.cpp similarity index 88% rename from samples/counter_collection/print_functional_counters.cpp rename to samples/counter_collection/print_functional_counters_client.cpp index 8621ff47b2..7986844c3f 100644 --- a/samples/counter_collection/print_functional_counters.cpp +++ b/samples/counter_collection/print_functional_counters_client.cpp @@ -22,7 +22,11 @@ #include "client.hpp" +#include +#include + #include +#include #include #include #include @@ -30,10 +34,6 @@ #include #include -#include -#include -#include - #define PRINT_ONLY_FAILING false /** @@ -245,7 +245,7 @@ get_agent_info() void dispatch_callback(rocprofiler_dispatch_counting_service_data_t dispatch_data, - rocprofiler_profile_config_id_t* config, + rocprofiler_counter_config_id_t* config, rocprofiler_user_data_t* /*user_data*/, void* /*callback_data_args*/) { @@ -282,40 +282,24 @@ dispatch_callback(rocprofiler_dispatch_counting_service_data_t dispatch_data, for(auto& found_counter : counters_needed) { - rocprofiler_counter_info_v0_t version; + rocprofiler_counter_info_v1_t info; - ROCPROFILER_CALL(rocprofiler_query_counter_info(found_counter, - ROCPROFILER_COUNTER_INFO_VERSION_0, - static_cast(&version)), - "Could not query counter_id"); - cap.expected_counter_names.emplace(found_counter.handle, std::string(version.name)); - size_t expected = 0; - ROCPROFILER_CALL(rocprofiler_query_counter_instance_count( - dispatch_data.dispatch_info.agent_id, found_counter, &expected), - "COULD NOT QUERY INSTANCES"); + ROCPROFILER_CALL( + rocprofiler_query_counter_info( + found_counter, ROCPROFILER_COUNTER_INFO_VERSION_1, static_cast(&info)), + "Could not query counter_id"); + cap.expected_counter_names.emplace(found_counter.handle, std::string(info.name)); cap.remaining.push_back(found_counter); - cap.expected.emplace(found_counter.handle, expected); + cap.expected.emplace(found_counter.handle, info.instance_ids_count); auto& info_vector = cap.expected_data_dims.emplace(found_counter.handle, validate_dim_presence{}) .first->second; - ROCPROFILER_CALL(rocprofiler_iterate_counter_dimensions( - found_counter, - [](rocprofiler_counter_id_t, - const rocprofiler_record_dimension_info_t* dim_info, - size_t num_dims, - void* user_data) { - validate_dim_presence* dim_presence = - static_cast(user_data); - for(size_t i = 0; i < num_dims; i++) - { - dim_presence->maybe_forward(dim_info[i]); - } - return ROCPROFILER_STATUS_SUCCESS; - }, - static_cast(&info_vector)), - "Could not fetch dimension info"); + for(uint64_t i = 0; i < info.dimensions_count; i++) + { + info_vector.maybe_forward(info.dimensions[i]); + } } if(cap.expected.empty()) { @@ -326,10 +310,10 @@ dispatch_callback(rocprofiler_dispatch_counting_service_data_t dispatch_data, } if(cap.remaining.empty()) return; - rocprofiler_profile_config_id_t profile = {.handle = 0}; + rocprofiler_counter_config_id_t profile = {.handle = 0}; // Select the next counter to collect. - if(rocprofiler_create_profile_config( + if(rocprofiler_create_counter_config( dispatch_data.dispatch_info.agent_id, &(cap.remaining.back()), 1, &profile) == ROCPROFILER_STATUS_SUCCESS) { @@ -362,7 +346,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_counting_service( + ROCPROFILER_CALL(rocprofiler_configure_buffer_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/samples/external_correlation_id_request/client.cpp b/samples/external_correlation_id_request/client.cpp index 2ec5002c4c..599105b5df 100644 --- a/samples/external_correlation_id_request/client.cpp +++ b/samples/external_correlation_id_request/client.cpp @@ -33,14 +33,6 @@ #include "client.hpp" -#include -#include -#include -#include -#include -#include -#include -#include #include #include diff --git a/samples/openmp_target/client.cpp b/samples/openmp_target/client.cpp index d081748a37..bd5e1ca80c 100644 --- a/samples/openmp_target/client.cpp +++ b/samples/openmp_target/client.cpp @@ -33,11 +33,6 @@ #include "client.hpp" -#include -#include -#include -#include -#include #include #include diff --git a/samples/pc_sampling/client.cpp b/samples/pc_sampling/client.cpp index 286f58562e..48637bdd24 100644 --- a/samples/pc_sampling/client.cpp +++ b/samples/pc_sampling/client.cpp @@ -34,9 +34,6 @@ #include "pcs.hpp" #include "utils.hpp" -#include -#include -#include #include #include diff --git a/source/docs/api-reference/counter_collection_services.rst b/source/docs/api-reference/counter_collection_services.rst index 33ec07b0a4..bbcbf473e5 100644 --- a/source/docs/api-reference/counter_collection_services.rst +++ b/source/docs/api-reference/counter_collection_services.rst @@ -65,7 +65,7 @@ After creating a context and buffer to store results in ``tool_init``, it is hig // 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_counting_service( + ROCPROFILER_CALL(rocprofiler_configure_buffer_dispatch_counting_service( ctx, buff, dispatch_callback, nullptr), "Could not setup buffered service"); @@ -145,21 +145,21 @@ Profile Setup { // Contains name and other attributes about the counter. // See API documentation for more info on the contents of this struct. - rocprofiler_counter_info_v0_t version; + rocprofiler_counter_info_v0_t info; ROCPROFILER_CALL( rocprofiler_query_counter_info( - counter, ROCPROFILER_COUNTER_INFO_VERSION_0, static_cast(&version)), + counter, ROCPROFILER_COUNTER_INFO_VERSION_0, static_cast(&info)), "Could not query info for counter"); } -4. After identifying the counters to be collected, construct a profile by passing a list of these counters to ``rocprofiler_create_profile_config``. +4. After identifying the counters to be collected, construct a profile by passing a list of these counters to ``rocprofiler_create_counter_config``. .. code-block:: cpp // Create and return the profile - rocprofiler_profile_config_id_t profile; - ROCPROFILER_CALL(rocprofiler_create_profile_config( + rocprofiler_counter_config_id_t profile; + ROCPROFILER_CALL(rocprofiler_create_counter_config( agent, counters_array, counters_array_count, &profile), "Could not construct profile cfg"); @@ -172,7 +172,7 @@ Profile Setup - Profile created is *only valid* for the agent it was created for. - Profiles are immutable. To collect a new counter set, construct a new profile. - A single profile can be used multiple times on the same agent. - - Counter Ids supplied to ``rocprofiler_create_profile_config`` are *agent-specific* and can't be used to construct profiles for other agents. + - Counter Ids supplied to ``rocprofiler_create_counter_config`` are *agent-specific* and can't be used to construct profiles for other agents. Dispatch Counting Callback -------------------------- @@ -183,7 +183,7 @@ When a kernel is dispatched, a dispatch callback is issued to the tool to allow void dispatch_callback(rocprofiler_dispatch_counting_service_data_t dispatch_data, - rocprofiler_profile_config_id_t* config, + rocprofiler_counter_config_id_t* config, rocprofiler_user_data_t* user_data, void* /*callback_data_args*/) @@ -197,9 +197,9 @@ This callback is invoked after the context starts and allows the tool to specify .. code-block:: cpp void - set_profile(rocprofiler_context_id_t context_id, - rocprofiler_agent_id_t agent, - rocprofiler_agent_set_profile_callback_t set_config, + set_profile(rocprofiler_context_id_t context_id, + rocprofiler_agent_id_t agent, + rocprofiler_device_counting_agent_cb_t set_config, void*) The profile to be used for this agent is specified by calling ``set_config(agent, profile)``. @@ -210,7 +210,7 @@ Buffered callback Data from collected counter values is returned through a buffered callback. The buffered callback routines are similar for dispatch and device counting except that some data such as kernel launch Ids is not available in device counting mode. Here is a sample iteration to print out counter collection data: .. code-block:: cpp - + for(size_t i = 0; i < num_headers; ++i) { auto* header = headers[i]; diff --git a/source/docs/rocprofiler-sdk.dox.in b/source/docs/rocprofiler-sdk.dox.in index 25a90eac69..4e73c1f3cc 100644 --- a/source/docs/rocprofiler-sdk.dox.in +++ b/source/docs/rocprofiler-sdk.dox.in @@ -124,7 +124,8 @@ WARN_LOGFILE = #--------------------------------------------------------------------------- INPUT = @SOURCE_DIR@/README.md \ @SOURCE_DIR@/source/include/rocprofiler-sdk \ - @SOURCE_DIR@/build-docs/source/include/rocprofiler-sdk/version.h + @SOURCE_DIR@/build-docs/source/include/rocprofiler-sdk/version.h \ + @SOURCE_DIR@/build-docs/source/include/rocprofiler-sdk/ext_version.h INPUT_ENCODING = UTF-8 INPUT_FILE_ENCODING = FILE_PATTERNS = *.h \ @@ -177,7 +178,7 @@ INPUT_FILTER = FILTER_PATTERNS = FILTER_SOURCE_FILES = NO FILTER_SOURCE_PATTERNS = -USE_MDFILE_AS_MAINPAGE = +USE_MDFILE_AS_MAINPAGE = FORTRAN_COMMENT_AFTER = 72 #--------------------------------------------------------------------------- # Configuration options related to source browsing @@ -202,10 +203,10 @@ IGNORE_PREFIX = GENERATE_HTML = YES HTML_OUTPUT = html HTML_FILE_EXTENSION = .html -HTML_HEADER = -HTML_FOOTER = -HTML_STYLESHEET = -HTML_EXTRA_STYLESHEET = +HTML_HEADER = +HTML_FOOTER = +HTML_STYLESHEET = +HTML_EXTRA_STYLESHEET = HTML_EXTRA_FILES = HTML_COLORSTYLE = LIGHT HTML_COLORSTYLE_HUE = 220 @@ -338,7 +339,7 @@ EXPAND_ONLY_PREDEF = NO SEARCH_INCLUDES = NO INCLUDE_PATH = @SOURCE_DIR@/source/include INCLUDE_FILE_PATTERNS = *.h \ - *.hpp + *.hpp PREDEFINED = "ROCPROFILER_API=" \ "ROCPROFILER_EXPORT=" \ "ROCPROFILER_IMPORT=" \ diff --git a/source/include/rocprofiler-sdk/CMakeLists.txt b/source/include/rocprofiler-sdk/CMakeLists.txt index 3f3f5bfeb7..a9184cbd22 100644 --- a/source/include/rocprofiler-sdk/CMakeLists.txt +++ b/source/include/rocprofiler-sdk/CMakeLists.txt @@ -6,18 +6,22 @@ configure_file(${CMAKE_CURRENT_LIST_DIR}/version.h.in ${CMAKE_CURRENT_BINARY_DIR}/version.h @ONLY) +configure_file(${CMAKE_CURRENT_LIST_DIR}/ext_version.h.in + ${CMAKE_CURRENT_BINARY_DIR}/ext_version.h @ONLY) + set(ROCPROFILER_HEADER_FILES # core headers rocprofiler.h # secondary headers agent.h - device_counting_service.h buffer.h buffer_tracing.h callback_tracing.h context.h + counter_config.h counters.h defines.h + device_counting_service.h dispatch_counting_service.h external_correlation.h fwd.h @@ -28,13 +32,13 @@ set(ROCPROFILER_HEADER_FILES marker.h ompt.h pc_sampling.h - profile_config.h - registration.h rccl.h + registration.h rocdecode.h rocjpeg.h spm.h - ${CMAKE_CURRENT_BINARY_DIR}/version.h) + ${CMAKE_CURRENT_BINARY_DIR}/version.h + ${CMAKE_CURRENT_BINARY_DIR}/ext_version.h) install( FILES ${ROCPROFILER_HEADER_FILES} @@ -52,3 +56,4 @@ add_subdirectory(cxx) add_subdirectory(kfd) add_subdirectory(amd_detail) add_subdirectory(experimental) +add_subdirectory(deprecated) diff --git a/source/include/rocprofiler-sdk/amd_detail/thread_trace_core.h b/source/include/rocprofiler-sdk/amd_detail/thread_trace_core.h index b77741a966..0852a73bef 100644 --- a/source/include/rocprofiler-sdk/amd_detail/thread_trace_core.h +++ b/source/include/rocprofiler-sdk/amd_detail/thread_trace_core.h @@ -36,7 +36,7 @@ ROCPROFILER_EXTERN_C_INIT * @{ */ -typedef enum +typedef enum rocprofiler_att_parameter_type_t { ROCPROFILER_ATT_PARAMETER_TARGET_CU = 0, ///< Select the Target CU or WGP ROCPROFILER_ATT_PARAMETER_SHADER_ENGINE_MASK, ///< Bitmask of shader engines. @@ -48,7 +48,7 @@ typedef enum ROCPROFILER_ATT_PARAMETER_LAST } rocprofiler_att_parameter_type_t; -typedef struct +typedef struct rocprofiler_att_parameter_t { rocprofiler_att_parameter_type_t type; union diff --git a/source/include/rocprofiler-sdk/amd_detail/thread_trace_dispatch.h b/source/include/rocprofiler-sdk/amd_detail/thread_trace_dispatch.h index 62fa3e4d7c..935e5f6fb0 100644 --- a/source/include/rocprofiler-sdk/amd_detail/thread_trace_dispatch.h +++ b/source/include/rocprofiler-sdk/amd_detail/thread_trace_dispatch.h @@ -37,7 +37,7 @@ ROCPROFILER_EXTERN_C_INIT * @{ */ -typedef enum +typedef enum rocprofiler_att_control_flags_t { ROCPROFILER_ATT_CONTROL_NONE = 0, ROCPROFILER_ATT_CONTROL_START_AND_STOP = 3 diff --git a/source/include/rocprofiler-sdk/buffer_tracing.h b/source/include/rocprofiler-sdk/buffer_tracing.h index 2bffc42c18..5f6ccde9ec 100644 --- a/source/include/rocprofiler-sdk/buffer_tracing.h +++ b/source/include/rocprofiler-sdk/buffer_tracing.h @@ -42,7 +42,7 @@ ROCPROFILER_EXTERN_C_INIT /** * @brief ROCProfiler Buffer HSA API Tracer Record. */ -typedef struct +typedef struct rocprofiler_buffer_tracing_hsa_api_record_t { uint64_t size; ///< size of this struct rocprofiler_buffer_tracing_kind_t kind; @@ -66,7 +66,7 @@ typedef struct /** * @brief ROCProfiler Buffer HIP API Tracer Record. */ -typedef struct +typedef struct rocprofiler_buffer_tracing_hip_api_record_t { uint64_t size; ///< size of this struct rocprofiler_buffer_tracing_kind_t kind; @@ -87,7 +87,7 @@ typedef struct /** * @brief ROCProfiler Buffer HIP API Tracer Record. */ -typedef struct +typedef struct rocprofiler_buffer_tracing_hip_api_ext_record_t { uint64_t size; ///< size of this struct rocprofiler_buffer_tracing_kind_t kind; @@ -167,7 +167,7 @@ typedef struct rocprofiler_buffer_tracing_ompt_record_t /** * @brief ROCProfiler Buffer Marker Tracer Record. */ -typedef struct +typedef struct rocprofiler_buffer_tracing_marker_api_record_t { uint64_t size; ///< size of this struct rocprofiler_buffer_tracing_kind_t kind; @@ -190,7 +190,7 @@ typedef struct /** * @brief ROCProfiler Buffer RCCL API Record. */ -typedef struct +typedef struct rocprofiler_buffer_tracing_rccl_api_record_t { uint64_t size; ///< size of this struct rocprofiler_buffer_tracing_kind_t kind; @@ -247,7 +247,7 @@ typedef struct rocprofiler_buffer_tracing_rocjpeg_api_record_t /** * @brief ROCProfiler Buffer Memory Copy Tracer Record. */ -typedef struct +typedef struct rocprofiler_buffer_tracing_memory_copy_record_t { uint64_t size; ///< size of this struct rocprofiler_buffer_tracing_kind_t kind; @@ -272,7 +272,7 @@ typedef struct /** * @brief ROCProfiler Buffer Memory Allocation Tracer Record. */ -typedef struct +typedef struct rocprofiler_buffer_tracing_memory_allocation_record_t { uint64_t size; ///< size of this struct rocprofiler_buffer_tracing_kind_t kind; @@ -328,7 +328,7 @@ typedef struct rocprofiler_buffer_tracing_page_migration_record_t /** * @brief ROCProfiler Buffer Scratch Memory Tracer Record */ -typedef struct +typedef struct rocprofiler_buffer_tracing_scratch_memory_record_t { uint64_t size; ///< size of this struct rocprofiler_buffer_tracing_kind_t kind; ///< ::ROCPROFILER_BUFFER_TRACING_SCRATCH_MEMORY @@ -345,7 +345,7 @@ typedef struct /** * @brief ROCProfiler Buffer Correlation ID Retirement Tracer Record. */ -typedef struct +typedef struct rocprofiler_buffer_tracing_correlation_id_retirement_record_t { uint64_t size; ///< size of this struct rocprofiler_buffer_tracing_kind_t kind; diff --git a/source/include/rocprofiler-sdk/callback_tracing.h b/source/include/rocprofiler-sdk/callback_tracing.h index 9c824871d5..fc03961da8 100644 --- a/source/include/rocprofiler-sdk/callback_tracing.h +++ b/source/include/rocprofiler-sdk/callback_tracing.h @@ -52,7 +52,7 @@ ROCPROFILER_EXTERN_C_INIT * @brief ROCProfiler Enumeration for code object storage types (identical values to * `hsa_ven_amd_loader_code_object_storage_type_t` enumeration) */ -typedef enum +typedef enum rocprofiler_code_object_storage_type_t { ROCPROFILER_CODE_OBJECT_STORAGE_TYPE_NONE = HSA_VEN_AMD_LOADER_CODE_OBJECT_STORAGE_TYPE_NONE, ROCPROFILER_CODE_OBJECT_STORAGE_TYPE_FILE = HSA_VEN_AMD_LOADER_CODE_OBJECT_STORAGE_TYPE_FILE, @@ -64,7 +64,7 @@ typedef enum /** * @brief ROCProfiler HSA API Callback Data. */ -typedef struct +typedef struct rocprofiler_callback_tracing_hsa_api_data_t { uint64_t size; ///< size of this struct rocprofiler_hsa_api_args_t args; @@ -74,7 +74,7 @@ typedef struct /** * @brief ROCProfiler HIP runtime and compiler API Tracer Callback Data. */ -typedef struct +typedef struct rocprofiler_callback_tracing_hip_api_data_t { uint64_t size; ///< size of this struct rocprofiler_hip_api_args_t args; @@ -89,7 +89,7 @@ typedef struct /** * @brief ROCProfiler OMPT Callback Data */ -typedef struct +typedef struct rocprofiler_callback_tracing_ompt_data_t { uint64_t size; ///< size of this struct rocprofiler_ompt_args_t args; @@ -98,7 +98,7 @@ typedef struct /** * @brief ROCProfiler Marker Tracer Callback Data. */ -typedef struct +typedef struct rocprofiler_callback_tracing_marker_api_data_t { uint64_t size; ///< size of this struct rocprofiler_marker_api_args_t args; @@ -108,7 +108,7 @@ typedef struct /** * @brief ROCProfiler RCCL API Callback Data. */ -typedef struct +typedef struct rocprofiler_callback_tracing_rccl_api_data_t { uint64_t size; ///< size of this struct rocprofiler_rccl_api_args_t args; @@ -138,7 +138,7 @@ typedef struct rocprofiler_callback_tracing_rocjpeg_api_data_t /** * @brief ROCProfiler Code Object Load Tracer Callback Record. */ -typedef struct +typedef struct rocprofiler_callback_tracing_code_object_load_data_t { uint64_t size; ///< size of this struct uint64_t code_object_id; ///< unique code object identifier @@ -245,7 +245,7 @@ typedef struct rocprofiler_callback_tracing_code_object_kernel_symbol_register_d } rocprofiler_callback_tracing_code_object_kernel_symbol_register_data_t; // rename struct -typedef struct +typedef struct rocprofiler_callback_tracing_code_object_host_kernel_symbol_register_data_t { uint64_t size; ///< size of this struct uint64_t host_function_id; ///< unique host function identifier value @@ -283,7 +283,7 @@ typedef struct rocprofiler_callback_tracing_kernel_dispatch_data_t * The timestamps in this record will only be non-zero in the ::ROCPROFILER_CALLBACK_PHASE_EXIT * callback */ -typedef struct +typedef struct rocprofiler_callback_tracing_memory_copy_data_t { uint64_t size; ///< size of this struct rocprofiler_timestamp_t start_timestamp; ///< start time in nanoseconds @@ -298,7 +298,7 @@ typedef struct /** * @brief ROCProfiler Memory Allocation Tracer Record. */ -typedef struct +typedef struct rocprofiler_callback_tracing_memory_allocation_data_t { uint64_t size; ///< size of this struct rocprofiler_timestamp_t start_timestamp; ///< start time in nanoseconds @@ -339,11 +339,11 @@ typedef struct rocprofiler_callback_tracing_runtime_initialization_data_t /** * @brief ROCProfiler Stream Handle Callback Data. */ -typedef struct +typedef struct rocprofiler_callback_tracing_hip_stream_data_t { uint64_t size; ///< size of this struct rocprofiler_stream_id_t stream_id; ///< HIP stream ID -} rocprofiler_callback_tracing_stream_handle_data_t; +} rocprofiler_callback_tracing_hip_stream_data_t; /** * @brief API Tracing callback function. This function is invoked twice per API function: once diff --git a/source/include/rocprofiler-sdk/profile_config.h b/source/include/rocprofiler-sdk/counter_config.h similarity index 66% rename from source/include/rocprofiler-sdk/profile_config.h rename to source/include/rocprofiler-sdk/counter_config.h index 5d6be8ebf2..f0a1c3cce4 100644 --- a/source/include/rocprofiler-sdk/profile_config.h +++ b/source/include/rocprofiler-sdk/counter_config.h @@ -29,51 +29,53 @@ ROCPROFILER_EXTERN_C_INIT /** - * @defgroup PROFILE_CONFIG Profile Configurations + * @defgroup COUNTER_CONFIG HW Counter Configurations * @brief Group one or more hardware counters into a unique handle * * @{ */ /** - * @brief Create Profile Configuration. A profile is bound to an agent but can - * be used across many contexts. The profile has a fixed set of counters + * @brief (experimental) Create Counter Configuration. A config is bound to an agent but can + * be used across many contexts. The config has a fixed set of counters * that are collected (and specified by counter_list). The available * counters for an agent can be queried using - * @ref rocprofiler_iterate_agent_supported_counters. An existing profile - * may be supplied via config_id to use as a base for the new profile. - * All counters in the existing profile will be copied over to the new - * profile. The existing profile will remain unmodified and usable with - * the new profile id being returned in config_id. + * @ref rocprofiler_iterate_agent_supported_counters. An existing config + * may be supplied via config_id to use as a base for the new config. + * All counters in the existing config will be copied over to the new + * config. The existing config will remain unmodified and usable with + * the new config id being returned in config_id. * * @param [in] agent_id Agent identifier * @param [in] counters_list List of GPU counters * @param [in] counters_count Size of counters list * @param [in,out] config_id Identifier for GPU counters group. If an existing - profile is supplied, that profiles counters will be copied - over to a new profile (returned via this id) + config is supplied, that profiles counters will be copied + over to a new config (returned via this id) * @return ::rocprofiler_status_t - * @retval ROCPROFILER_STATUS_SUCCESS if profile created - * @retval ROCPROFILER_STATUS_ERROR if profile could not be created + * @retval ROCPROFILER_STATUS_SUCCESS if config created + * @retval ROCPROFILER_STATUS_ERROR if config could not be created * */ +ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_status_t -rocprofiler_create_profile_config(rocprofiler_agent_id_t agent_id, +rocprofiler_create_counter_config(rocprofiler_agent_id_t agent_id, rocprofiler_counter_id_t* counters_list, size_t counters_count, - rocprofiler_profile_config_id_t* config_id) ROCPROFILER_API + rocprofiler_counter_config_id_t* config_id) ROCPROFILER_API ROCPROFILER_NONNULL(4); /** - * @brief Destroy Profile Configuration. + * @brief (experimental) Destroy Profile Configuration. * * @param [in] config_id * @return ::rocprofiler_status_t - * @retval ROCPROFILER_STATUS_SUCCESS if profile destroyed - * @retval ROCPROFILER_STATUS_ERROR if profile could not be destroyed + * @retval ROCPROFILER_STATUS_SUCCESS if config destroyed + * @retval ROCPROFILER_STATUS_ERROR if config could not be destroyed */ +ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_status_t -rocprofiler_destroy_profile_config(rocprofiler_profile_config_id_t config_id) ROCPROFILER_API; +rocprofiler_destroy_counter_config(rocprofiler_counter_config_id_t config_id) ROCPROFILER_API; /** @} */ diff --git a/source/include/rocprofiler-sdk/counters.h b/source/include/rocprofiler-sdk/counters.h index ce6b68f7fa..cea7050b98 100644 --- a/source/include/rocprofiler-sdk/counters.h +++ b/source/include/rocprofiler-sdk/counters.h @@ -26,6 +26,8 @@ #include #include +#include + ROCPROFILER_EXTERN_C_INIT /** @@ -35,20 +37,66 @@ ROCPROFILER_EXTERN_C_INIT */ /** - * @brief Query counter id information from record_id. + * @brief (experimental) Counter info struct version 0 + */ +typedef struct ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_counter_info_v0_t +{ + rocprofiler_counter_id_t id; ///< Id of this counter + const char* name; ///< Name of the counter + const char* description; ///< Description of the counter + const char* block; ///< Block of the counter (non-derived only) + const char* expression; ///< Counter expression (derived counters only) + uint8_t is_constant : 1; ///< If this counter is HW constant + uint8_t is_derived : 1; ///< If this counter is a derived counter +} rocprofiler_counter_info_v0_t; + +/** + * @brief (experimental) Counter info struct version 1. Combines information from + * ::rocprofiler_counter_info_v0_t with the dimension information. + */ +typedef struct ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_counter_info_v1_t +{ + rocprofiler_counter_id_t id; ///< Id of this counter + const char* name; ///< Name of the counter + const char* description; ///< Description of the counter + const char* block; ///< Block of the counter (non-derived only) + const char* expression; ///< Counter expression (derived counters only) + uint8_t is_constant : 1; ///< If this counter is HW constant + uint8_t is_derived : 1; ///< If this counter is a derived counter + + uint64_t dimensions_count; + const rocprofiler_counter_record_dimension_info_t* dimensions; + uint64_t instance_ids_count; + const rocprofiler_counter_instance_id_t* instance_ids; + + /// @var dimensions_count + /// @brief Number of dimensions for the counter + /// + /// @var dimensions + /// @brief Dimension information of the counter + /// + /// @var instance_ids_count + /// @brief Number of instance ids for the counter + /// + /// @var instance_ids + /// @brief Instance ids that can be generated by the counter +} rocprofiler_counter_info_v1_t; + +/** + * @brief (experimental) Query counter id information from record_id. * * @param [in] id record id from rocprofiler_record_counter_t * @param [out] counter_id counter id associated with the record * @return ::rocprofiler_status_t * @retval ROCPROFILER_STATUS_SUCCESS if id decoded */ -rocprofiler_status_t +ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_status_t rocprofiler_query_record_counter_id(rocprofiler_counter_instance_id_t id, rocprofiler_counter_id_t* counter_id) ROCPROFILER_API ROCPROFILER_NONNULL(2); /** - * @brief Query dimension position from record_id. If the dimension does not exist + * @brief (experimental) Query dimension position from record_id. If the dimension does not exist * in the counter, the return will be 0. * * @param [in] id record id from @ref rocprofiler_record_counter_t @@ -58,46 +106,13 @@ rocprofiler_query_record_counter_id(rocprofiler_counter_instance_id_t id, * @return ::rocprofiler_status_t * @retval ROCPROFILER_STATUS_SUCCESS if dimension decoded */ -rocprofiler_status_t +ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_status_t rocprofiler_query_record_dimension_position(rocprofiler_counter_instance_id_t id, rocprofiler_counter_dimension_id_t dim, size_t* pos) ROCPROFILER_API ROCPROFILER_NONNULL(3); /** - * @brief Callback that gives a list of available dimensions for a counter - * - * @param [in] id Counter id the dimension data is for - * @param [in] dim_info An array of dimensions for the counter - * @ref rocprofiler_iterate_counter_dimensions was called on. - * @param [in] num_dims Number of dimensions - * @param [in] user_data User data supplied by - * @ref rocprofiler_iterate_agent_supported_counters - */ -typedef rocprofiler_status_t (*rocprofiler_available_dimensions_cb_t)( - rocprofiler_counter_id_t id, - const rocprofiler_record_dimension_info_t* dim_info, - size_t num_dims, - void* user_data); - -/** - * @brief Return information about the dimensions that exists for a specific counter - * and the extent of each dimension. - * - * @param [in] id counter id to query dimension info for. - * @param [in] info_cb Callback to return dimension information for counter - * @param [in] user_data data to pass into the callback - * @return ::rocprofiler_status_t - * @retval ROCPROFILER_STATUS_SUCCESS if dimension exists - * @retval ROCPROFILER_STATUS_ERROR_COUNTER_NOT_FOUND if counter is not found - * @retval ROCPROFILER_STATUS_ERROR_DIM_NOT_FOUND if counter does not have this dimension - */ -rocprofiler_status_t -rocprofiler_iterate_counter_dimensions(rocprofiler_counter_id_t id, - rocprofiler_available_dimensions_cb_t info_cb, - void* user_data) ROCPROFILER_API; - -/** - * @brief Query Counter info such as name or description. + * @brief (experimental) Query Counter info such as name or description. * * @param [in] counter_id counter to get info for * @param [in] version Version of struct in info, see @ref rocprofiler_counter_info_version_id_t for @@ -108,29 +123,13 @@ rocprofiler_iterate_counter_dimensions(rocprofiler_counter_id_t id, * @retval ROCPROFILER_STATUS_ERROR_COUNTER_NOT_FOUND if counter not found * @retval ROCPROFILER_STATUS_ERROR_INCOMPATIBLE_ABI Version is not supported */ -rocprofiler_status_t +ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_status_t rocprofiler_query_counter_info(rocprofiler_counter_id_t counter_id, rocprofiler_counter_info_version_id_t version, void* info) ROCPROFILER_API ROCPROFILER_NONNULL(3); /** - * @brief This call returns the number of instances specific counter contains. - * - * @param [in] agent_id rocprofiler agent identifier - * @param [in] counter_id counter id (obtained from iterate_agent_supported_counters) - * @param [out] instance_count number of instances the counter has - * @return ::rocprofiler_status_t - * @retval ROCPROFILER_STATUS_SUCCESS if counter found - * @retval ROCPROFILER_STATUS_ERROR_COUNTER_NOT_FOUND if counter not found - */ -rocprofiler_status_t -rocprofiler_query_counter_instance_count(rocprofiler_agent_id_t agent_id, - rocprofiler_counter_id_t counter_id, - size_t* instance_count) ROCPROFILER_API - ROCPROFILER_NONNULL(3); - -/** - * @brief Callback that gives a list of counters available on an agent. The + * @brief (experimental) Callback that gives a list of counters available on an agent. The * counters variable is owned by rocprofiler and should not be free'd. * * @param [in] agent_id Agent ID of the current callback @@ -140,14 +139,14 @@ rocprofiler_query_counter_instance_count(rocprofiler_agent_id_t agent_id, * @param [in] user_data User data supplied by * @ref rocprofiler_iterate_agent_supported_counters */ -typedef rocprofiler_status_t (*rocprofiler_available_counters_cb_t)( +ROCPROFILER_SDK_EXPERIMENTAL typedef rocprofiler_status_t (*rocprofiler_available_counters_cb_t)( rocprofiler_agent_id_t agent_id, rocprofiler_counter_id_t* counters, size_t num_counters, void* user_data); /** - * @brief Query Agent Counters Availability. + * @brief (experimental) Query Agent Counters Availability. * * @param [in] agent_id GPU agent identifier * @param [in] cb callback to caller to get counters @@ -156,20 +155,17 @@ typedef rocprofiler_status_t (*rocprofiler_available_counters_cb_t)( * @retval ROCPROFILER_STATUS_SUCCESS if counters found for agent * @retval ROCPROFILER_STATUS_ERROR if no counters found for agent */ -rocprofiler_status_t +ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_status_t rocprofiler_iterate_agent_supported_counters(rocprofiler_agent_id_t agent_id, rocprofiler_available_counters_cb_t cb, void* user_data) ROCPROFILER_API ROCPROFILER_NONNULL(2); -/** @} */ - /** - * @brief Creates a new counter based on a derived metric provided. The counter will only - * be available for counter collection profiles created after the addition of this counter. - * Due to the regeneration of internal ASTs and dimension cache, this call may be slow and - * should generally be avoided in performance sensitive code blocks (i.e. dispatch - * callbacks). + * @brief (experimental) Creates a new counter based on a derived metric provided. The counter will + * only be available for counter collection profiles created after the addition of this counter. Due + * to the regeneration of internal ASTs and dimension cache, this call may be slow and should + * generally be avoided in performance sensitive code blocks (i.e. dispatch callbacks). * * @param [in] name The name of the new counter. * @param [in] name_len The length of the counter name. @@ -186,7 +182,7 @@ rocprofiler_iterate_agent_supported_counters(rocprofiler_agent_id_t * @retval ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT if a counter argument is incorrect * @retval ROCPROFILER_STATUS_ERROR_AGENT_NOT_FOUND if the agent is not found */ -rocprofiler_status_t +ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_status_t rocprofiler_create_counter(const char* name, size_t name_len, const char* expr, @@ -196,4 +192,6 @@ rocprofiler_create_counter(const char* name, rocprofiler_agent_id_t agent, rocprofiler_counter_id_t* counter_id) ROCPROFILER_NONNULL(1, 3, 8); +/** @} */ + ROCPROFILER_EXTERN_C_FINI diff --git a/source/include/rocprofiler-sdk/cxx/CMakeLists.txt b/source/include/rocprofiler-sdk/cxx/CMakeLists.txt index bacc4678a2..fb087865e2 100644 --- a/source/include/rocprofiler-sdk/cxx/CMakeLists.txt +++ b/source/include/rocprofiler-sdk/cxx/CMakeLists.txt @@ -4,7 +4,7 @@ # # set(ROCPROFILER_CXX_HEADER_FILES hash.hpp name_info.hpp operators.hpp perfetto.hpp - utility.hpp serialization.hpp) + utility.hpp serialization.hpp version.hpp) install( FILES ${ROCPROFILER_CXX_HEADER_FILES} diff --git a/source/include/rocprofiler-sdk/cxx/hash.hpp b/source/include/rocprofiler-sdk/cxx/hash.hpp index 6fd1ea1868..6ba6eade8f 100644 --- a/source/include/rocprofiler-sdk/cxx/hash.hpp +++ b/source/include/rocprofiler-sdk/cxx/hash.hpp @@ -63,7 +63,7 @@ ROCPROFILER_CXX_SPECIALIZE_HANDLE_HASHER(rocprofiler_queue_id_t) ROCPROFILER_CXX_SPECIALIZE_HANDLE_HASHER(rocprofiler_stream_id_t) ROCPROFILER_CXX_SPECIALIZE_HANDLE_HASHER(rocprofiler_buffer_id_t) ROCPROFILER_CXX_SPECIALIZE_HANDLE_HASHER(rocprofiler_counter_id_t) -ROCPROFILER_CXX_SPECIALIZE_HANDLE_HASHER(rocprofiler_profile_config_id_t) +ROCPROFILER_CXX_SPECIALIZE_HANDLE_HASHER(rocprofiler_counter_config_id_t) ROCPROFILER_CXX_SPECIALIZE_HANDLE_HASHER(rocprofiler_callback_thread_t) ROCPROFILER_CXX_SPECIALIZE_HANDLE_HASHER(hsa_agent_t) ROCPROFILER_CXX_SPECIALIZE_HANDLE_HASHER(hsa_signal_t) diff --git a/source/include/rocprofiler-sdk/cxx/operators.hpp b/source/include/rocprofiler-sdk/cxx/operators.hpp index 67ee89c11d..f6a91f8b8d 100644 --- a/source/include/rocprofiler-sdk/cxx/operators.hpp +++ b/source/include/rocprofiler-sdk/cxx/operators.hpp @@ -28,6 +28,7 @@ #include #include #include +#include #include @@ -100,7 +101,7 @@ ROCPROFILER_CXX_DECLARE_OPERATORS(rocprofiler_queue_id_t) ROCPROFILER_CXX_DECLARE_OPERATORS(rocprofiler_stream_id_t) ROCPROFILER_CXX_DECLARE_OPERATORS(rocprofiler_buffer_id_t) ROCPROFILER_CXX_DECLARE_OPERATORS(rocprofiler_counter_id_t) -ROCPROFILER_CXX_DECLARE_OPERATORS(rocprofiler_profile_config_id_t) +ROCPROFILER_CXX_DECLARE_OPERATORS(rocprofiler_counter_config_id_t) ROCPROFILER_CXX_DECLARE_OPERATORS(rocprofiler_callback_thread_t) ROCPROFILER_CXX_DECLARE_OPERATORS(hsa_agent_t) ROCPROFILER_CXX_DECLARE_OPERATORS(hsa_signal_t) @@ -109,6 +110,8 @@ ROCPROFILER_CXX_DECLARE_OPERATORS(const rocprofiler_agent_v0_t&) ROCPROFILER_CXX_DECLARE_OPERATORS(rocprofiler_dim3_t) ROCPROFILER_CXX_DECLARE_OPERATORS(hsa_region_t) ROCPROFILER_CXX_DECLARE_OPERATORS(hsa_amd_memory_pool_t) +ROCPROFILER_CXX_DECLARE_OPERATORS(const rocprofiler_counter_record_dimension_info_t&) +ROCPROFILER_CXX_DECLARE_OPERATORS(rocprofiler_version_triplet_t) // definitions of operator== ROCPROFILER_CXX_DEFINE_EQ_HANDLE_OPERATOR(rocprofiler_context_id_t) @@ -118,7 +121,7 @@ ROCPROFILER_CXX_DEFINE_EQ_HANDLE_OPERATOR(rocprofiler_queue_id_t) ROCPROFILER_CXX_DEFINE_EQ_HANDLE_OPERATOR(rocprofiler_stream_id_t) ROCPROFILER_CXX_DEFINE_EQ_HANDLE_OPERATOR(rocprofiler_buffer_id_t) ROCPROFILER_CXX_DEFINE_EQ_HANDLE_OPERATOR(rocprofiler_counter_id_t) -ROCPROFILER_CXX_DEFINE_EQ_HANDLE_OPERATOR(rocprofiler_profile_config_id_t) +ROCPROFILER_CXX_DEFINE_EQ_HANDLE_OPERATOR(rocprofiler_counter_config_id_t) ROCPROFILER_CXX_DEFINE_EQ_HANDLE_OPERATOR(rocprofiler_callback_thread_t) ROCPROFILER_CXX_DEFINE_EQ_HANDLE_OPERATOR(hsa_agent_t) ROCPROFILER_CXX_DEFINE_EQ_HANDLE_OPERATOR(hsa_signal_t) @@ -138,6 +141,20 @@ operator==(rocprofiler_dim3_t lhs, rocprofiler_dim3_t rhs) return std::tie(lhs.x, lhs.y, lhs.z) == std::tie(rhs.x, rhs.y, rhs.z); } +inline bool +operator==(rocprofiler_counter_record_dimension_info_t lhs, + rocprofiler_counter_record_dimension_info_t rhs) +{ + return std::tie(lhs.id, lhs.instance_size, lhs.name) == + std::tie(rhs.id, rhs.instance_size, rhs.name); +} + +inline bool +operator==(rocprofiler_version_triplet_t lhs, rocprofiler_version_triplet_t rhs) +{ + return std::tie(lhs.major, lhs.minor, lhs.patch) == std::tie(rhs.major, rhs.minor, rhs.patch); +} + // definitions of operator!= ROCPROFILER_CXX_DEFINE_NE_OPERATOR(rocprofiler_context_id_t) ROCPROFILER_CXX_DEFINE_NE_OPERATOR(rocprofiler_address_t) @@ -146,7 +163,7 @@ ROCPROFILER_CXX_DEFINE_NE_OPERATOR(rocprofiler_queue_id_t) ROCPROFILER_CXX_DEFINE_NE_OPERATOR(rocprofiler_stream_id_t) ROCPROFILER_CXX_DEFINE_NE_OPERATOR(rocprofiler_buffer_id_t) ROCPROFILER_CXX_DEFINE_NE_OPERATOR(rocprofiler_counter_id_t) -ROCPROFILER_CXX_DEFINE_NE_OPERATOR(rocprofiler_profile_config_id_t) +ROCPROFILER_CXX_DEFINE_NE_OPERATOR(rocprofiler_counter_config_id_t) ROCPROFILER_CXX_DEFINE_NE_OPERATOR(rocprofiler_callback_thread_t) ROCPROFILER_CXX_DEFINE_NE_OPERATOR(hsa_agent_t) ROCPROFILER_CXX_DEFINE_NE_OPERATOR(hsa_signal_t) @@ -155,6 +172,7 @@ ROCPROFILER_CXX_DEFINE_NE_OPERATOR(const rocprofiler_agent_v0_t&) ROCPROFILER_CXX_DEFINE_NE_OPERATOR(rocprofiler_dim3_t) ROCPROFILER_CXX_DEFINE_NE_OPERATOR(hsa_region_t) ROCPROFILER_CXX_DEFINE_NE_OPERATOR(hsa_amd_memory_pool_t) +ROCPROFILER_CXX_DEFINE_NE_OPERATOR(rocprofiler_version_triplet_t) // definitions of operator< ROCPROFILER_CXX_DEFINE_LT_HANDLE_OPERATOR(rocprofiler_context_id_t) @@ -164,7 +182,7 @@ ROCPROFILER_CXX_DEFINE_LT_HANDLE_OPERATOR(rocprofiler_queue_id_t) ROCPROFILER_CXX_DEFINE_LT_HANDLE_OPERATOR(rocprofiler_stream_id_t) ROCPROFILER_CXX_DEFINE_LT_HANDLE_OPERATOR(rocprofiler_buffer_id_t) ROCPROFILER_CXX_DEFINE_LT_HANDLE_OPERATOR(rocprofiler_counter_id_t) -ROCPROFILER_CXX_DEFINE_LT_HANDLE_OPERATOR(rocprofiler_profile_config_id_t) +ROCPROFILER_CXX_DEFINE_LT_HANDLE_OPERATOR(rocprofiler_counter_config_id_t) ROCPROFILER_CXX_DEFINE_LT_HANDLE_OPERATOR(rocprofiler_callback_thread_t) ROCPROFILER_CXX_DEFINE_LT_HANDLE_OPERATOR(hsa_agent_t) ROCPROFILER_CXX_DEFINE_LT_HANDLE_OPERATOR(hsa_signal_t) @@ -172,6 +190,14 @@ ROCPROFILER_CXX_DEFINE_LT_HANDLE_OPERATOR(hsa_executable_t) ROCPROFILER_CXX_DEFINE_LT_HANDLE_OPERATOR(hsa_region_t) ROCPROFILER_CXX_DEFINE_LT_HANDLE_OPERATOR(hsa_amd_memory_pool_t) +inline bool +operator<(const rocprofiler_counter_record_dimension_info_t& lhs, + const rocprofiler_counter_record_dimension_info_t& rhs) +{ + return std::tie(lhs.id, lhs.instance_size, lhs.name) < + std::tie(rhs.id, rhs.instance_size, rhs.name); +} + inline bool operator<(const rocprofiler_agent_v0_t& lhs, const rocprofiler_agent_v0_t& rhs) { @@ -189,6 +215,12 @@ operator<(rocprofiler_dim3_t lhs, rocprofiler_dim3_t rhs) : (lhs_m < rhs_m); } +inline bool +operator<(rocprofiler_version_triplet_t lhs, rocprofiler_version_triplet_t rhs) +{ + return std::tie(lhs.major, lhs.minor, lhs.patch) < std::tie(rhs.major, rhs.minor, rhs.patch); +} + // definitions of operator>, operator<=, operator>= ROCPROFILER_CXX_DEFINE_COMPARE_OPERATORS(rocprofiler_context_id_t) ROCPROFILER_CXX_DEFINE_COMPARE_OPERATORS(rocprofiler_address_t) @@ -197,7 +229,7 @@ ROCPROFILER_CXX_DEFINE_COMPARE_OPERATORS(rocprofiler_queue_id_t) ROCPROFILER_CXX_DEFINE_COMPARE_OPERATORS(rocprofiler_stream_id_t) ROCPROFILER_CXX_DEFINE_COMPARE_OPERATORS(rocprofiler_buffer_id_t) ROCPROFILER_CXX_DEFINE_COMPARE_OPERATORS(rocprofiler_counter_id_t) -ROCPROFILER_CXX_DEFINE_COMPARE_OPERATORS(rocprofiler_profile_config_id_t) +ROCPROFILER_CXX_DEFINE_COMPARE_OPERATORS(rocprofiler_counter_config_id_t) ROCPROFILER_CXX_DEFINE_COMPARE_OPERATORS(rocprofiler_callback_thread_t) ROCPROFILER_CXX_DEFINE_COMPARE_OPERATORS(hsa_agent_t) ROCPROFILER_CXX_DEFINE_COMPARE_OPERATORS(hsa_signal_t) @@ -206,6 +238,7 @@ ROCPROFILER_CXX_DEFINE_COMPARE_OPERATORS(const rocprofiler_agent_v0_t&) ROCPROFILER_CXX_DEFINE_COMPARE_OPERATORS(rocprofiler_dim3_t) ROCPROFILER_CXX_DEFINE_COMPARE_OPERATORS(hsa_region_t) ROCPROFILER_CXX_DEFINE_COMPARE_OPERATORS(hsa_amd_memory_pool_t) +ROCPROFILER_CXX_DEFINE_COMPARE_OPERATORS(rocprofiler_version_triplet_t) // cleanup defines #undef ROCPROFILER_CXX_DECLARE_OPERATORS diff --git a/source/include/rocprofiler-sdk/cxx/serialization.hpp b/source/include/rocprofiler-sdk/cxx/serialization.hpp index 398c71a570..d385f34a58 100644 --- a/source/include/rocprofiler-sdk/cxx/serialization.hpp +++ b/source/include/rocprofiler-sdk/cxx/serialization.hpp @@ -1235,6 +1235,30 @@ save(ArchiveT& ar, rocprofiler_counter_info_v0_t data) ROCP_SDK_SAVE_DATA_CSTR(expression); } +template +void +save(ArchiveT& ar, rocprofiler_counter_info_v1_t data) +{ + ROCP_SDK_SAVE_DATA_FIELD(id); + ROCP_SDK_SAVE_DATA_BITFIELD("is_constant", is_constant); + ROCP_SDK_SAVE_DATA_BITFIELD("is_derived", is_derived); + ROCP_SDK_SAVE_DATA_CSTR(name); + ROCP_SDK_SAVE_DATA_CSTR(description); + ROCP_SDK_SAVE_DATA_CSTR(block); + ROCP_SDK_SAVE_DATA_CSTR(expression); + + auto convert = [](const auto* val, uint64_t sz) { + using data_type = std::remove_cv_t>; + auto retval = std::vector{}; + for(uint64_t i = 0; i < sz; ++i) + retval.emplace_back(val[i]); + return retval; + }; + + ROCP_SDK_SAVE_VALUE("dims", convert(data.dimensions, data.dimensions_count)); + ROCP_SDK_SAVE_VALUE("instances", convert(data.instance_ids, data.instance_ids_count)); +} + template void save(ArchiveT& ar, rocprofiler_record_dimension_info_t data) diff --git a/source/include/rocprofiler-sdk/cxx/version.hpp b/source/include/rocprofiler-sdk/cxx/version.hpp new file mode 100644 index 0000000000..839bfbc181 --- /dev/null +++ b/source/include/rocprofiler-sdk/cxx/version.hpp @@ -0,0 +1,59 @@ +// MIT License +// +// Copyright (c) 2023-2025 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 + +namespace rocprofiler +{ +namespace sdk +{ +namespace version +{ +/** + * @brief This function extracts the `..` from a single integer which encodes + * this triplet by a factor of `N`. E.g. version 3.2.1 using N=100 is represented by the value of + * 30201; version 3.2.1 using N=1000 is represented by the value of 3002001. + * + * For a given factor `N`, the major version can be extracted by + * dividing the version value by (N * N); the minor version can be extracted via + * computing the modulus of (N * N) and then divided by N; the patch version is simply the modulus + * of N. + */ +template +constexpr rocprofiler_version_triplet_t +compute_version_triplet(uint64_t version) +{ + constexpr auto factor = FactorV; + + return rocprofiler_version_triplet_t{ + .major = static_cast(version / (factor * factor)), + .minor = static_cast((version % (factor * factor)) / factor), + .patch = static_cast(version % factor)}; +} +} // namespace version +} // namespace sdk +} // namespace rocprofiler diff --git a/source/include/rocprofiler-sdk/defines.h b/source/include/rocprofiler-sdk/defines.h index a7c6441d42..142eecea74 100644 --- a/source/include/rocprofiler-sdk/defines.h +++ b/source/include/rocprofiler-sdk/defines.h @@ -48,6 +48,40 @@ /** @} */ +/** + * @def ROCPROFILER_SDK_BETA_COMPAT + * @brief rocprofiler-sdk clients (i.e. tool using rocprofiler-sdk) should set this definition to 1 + * before including any rocprofiler-sdk header if it wants rocprofiler-sdk to provide preprocessor + * definitions to help with compilation support for tools prior to v1.0.0 release. Note: for v1.0.0 + * release, rocprofiler-sdk sets the ppdef to 1 by default. Eventually, rocprofiler-sdk will remove + * defining this value and it will be up to the tools to define this value. + * + * For example in version 0.6.0, there was a function `rocprofiler_create_profile_config` and, prior + * to the 1.0.0 release, this function was renamed to `rocprofiler_create_counter_config`. + * @addtogroup VERSIONING_GROUP + * + * @def ROCPROFILER_SDK_BETA_COMPAT_SUPPORTED + * @brief rocprofiler-sdk will set this preprocessor definition to 1 if it can honor + * ::ROCPROFILER_SDK_BETA_COMPAT set to 1. Once backwards compatibility with the beta + * rocprofiler-sdk can no longer be supported, this will always be set to 0. + * @addtogroup VERSIONING_GROUP + * + * @def ROCPROFILER_SDK_DEPRECATED_WARNINGS + * @brief Set this preprocessor definition to 0 to silent compiler warnings when using features that + * are marked as deprecated. By default, rocprofiler-sdk defines this to equal to 1. + * @addtogroup VERSIONING_GROUP + * + * @def ROCPROFILER_SDK_EXPERIMENTAL_WARNINGS + * @brief Set this preprocessor definition to 1 to enable compiler warnings when using experimental + * features. @see ::ROCPROFILER_SDK_EXPERIMENTAL + * @addtogroup VERSIONING_GROUP + * + * @def ROCPROFILER_SDK_EXPERIMENTAL + * @brief When this attribute is added to a type, object, expression, etc., the developer should be + * aware that the API and/or ABI is subject to change in subsequent releases. + * @addtogroup VERSIONING_GROUP + */ + #if !defined(ROCPROFILER_ATTRIBUTE) # if defined(_MSC_VER) # define ROCPROFILER_ATTRIBUTE(...) __declspec(__VA_ARGS__) @@ -137,3 +171,47 @@ #else # define ROCPROFILER_UINT64_C(value) UINT64_C(value) #endif + +#if defined(__cplusplus) && __cplusplus >= 201402L +# define ROCPROFILER_SDK_DEPRECATED_MESSAGE(...) [[deprecated(__VA_ARGS__)]] +#elif !defined(__cplusplus) && defined(__STDC_VERSION__) && __STDC_VERSION__ >= 202311L +# define ROCPROFILER_SDK_DEPRECATED_MESSAGE(...) [[deprecated(__VA_ARGS__)]] +#else +# define ROCPROFILER_SDK_DEPRECATED_MESSAGE(...) ROCPROFILER_ATTRIBUTE(deprecated) +#endif + +// TODO(jomadsen): uncomment below code before v1.0.0 +// #if !defined(ROCPROFILER_SDK_DEPRECATED_WARNINGS) +// # define ROCPROFILER_SDK_DEPRECATED_WARNINGS 1 +// #endif + +#if defined(ROCPROFILER_SDK_DEPRECATED_WARNINGS) && ROCPROFILER_SDK_DEPRECATED_WARNINGS > 0 +# define ROCPROFILER_SDK_DEPRECATED(...) ROCPROFILER_SDK_DEPRECATED_MESSAGE(__VA_ARGS__) +#else +# define ROCPROFILER_SDK_DEPRECATED(...) +#endif + +#define ROCPROFILER_SDK_EXPERIMENTAL_MESSAGE \ + ROCPROFILER_SDK_DEPRECATED_MESSAGE( \ + "Note: this feature has been marked as experimental. Define " \ + "ROCPROFILER_SDK_EXPERIMENTAL_WARNINGS=0 to silence this message.") + +#if defined(ROCPROFILER_SDK_EXPERIMENTAL_WARNINGS) && ROCPROFILER_SDK_EXPERIMENTAL_WARNINGS > 0 +# define ROCPROFILER_SDK_EXPERIMENTAL ROCPROFILER_SDK_EXPERIMENTAL_MESSAGE +#else +# define ROCPROFILER_SDK_EXPERIMENTAL +#endif + +// +// if ROCPROFILER_SDK_BETA_COMPAT is > 0, provide some macros to help with compatibility. +// For 1.0.0 release, we define this by default +// +#if !defined(ROCPROFILER_SDK_BETA_COMPAT) +# define ROCPROFILER_SDK_BETA_COMPAT 1 +#endif + +// rocprofiler-sdk will set ROCPROFILER_SDK_BETA_COMPAT_SUPPORTED to 1 if it can support +// compatibility with rocprofiler-sdk < v1.0.0 +#if defined(ROCPROFILER_SDK_BETA_COMPAT) && ROCPROFILER_SDK_BETA_COMPAT > 0 +# define ROCPROFILER_SDK_BETA_COMPAT_SUPPORTED 1 +#endif diff --git a/source/include/rocprofiler-sdk/deprecated/CMakeLists.txt b/source/include/rocprofiler-sdk/deprecated/CMakeLists.txt new file mode 100644 index 0000000000..ecb0865c13 --- /dev/null +++ b/source/include/rocprofiler-sdk/deprecated/CMakeLists.txt @@ -0,0 +1,9 @@ +# +# headers containing deprecated functions, types, etc. +# +set(ROCPROFILER_DEPRECATED_HEADER_FILES counters.h profile_config.h) + +install( + FILES ${ROCPROFILER_DEPRECATED_HEADER_FILES} + DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}/rocprofiler-sdk/deprecated + COMPONENT development) diff --git a/source/include/rocprofiler-sdk/deprecated/counters.h b/source/include/rocprofiler-sdk/deprecated/counters.h new file mode 100644 index 0000000000..a95ca04a3c --- /dev/null +++ b/source/include/rocprofiler-sdk/deprecated/counters.h @@ -0,0 +1,87 @@ +// MIT License +// +// Copyright (c) 2023-2025 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 +#include + +ROCPROFILER_EXTERN_C_INIT + +/** + * @brief (deprecated) Callback that gives a list of available dimensions for a counter + * + * @param [in] id Counter id the dimension data is for + * @param [in] dim_info An array of dimensions for the counter + * ::rocprofiler_iterate_counter_dimensions was called on. + * @param [in] num_dims Number of dimensions + * @param [in] user_data User data supplied by + * ::rocprofiler_iterate_agent_supported_counters + */ +ROCPROFILER_SDK_DEPRECATED( + "Function using this alias has been deprecated. See rocprofiler_iterate_counter_dimensions") +typedef rocprofiler_status_t (*rocprofiler_available_dimensions_cb_t)( + rocprofiler_counter_id_t id, + const rocprofiler_record_dimension_info_t* dim_info, + size_t num_dims, + void* user_data); + +/** + * @brief (deprecated) Return information about the dimensions that exists for a specific counter + * and the extent of each dimension. + * + * @param [in] id counter id to query dimension info for. + * @param [in] info_cb Callback to return dimension information for counter + * @param [in] user_data data to pass into the callback + * @return ::rocprofiler_status_t + * @retval ROCPROFILER_STATUS_SUCCESS if dimension exists + * @retval ROCPROFILER_STATUS_ERROR_COUNTER_NOT_FOUND if counter is not found + * @retval ROCPROFILER_STATUS_ERROR_DIM_NOT_FOUND if counter does not have this dimension + */ +ROCPROFILER_SDK_DEPRECATED("Information now available in rocprofiler_counter_info_v1_t. " + "This function will be removed in the future.") +rocprofiler_status_t +rocprofiler_iterate_counter_dimensions(rocprofiler_counter_id_t id, + rocprofiler_available_dimensions_cb_t info_cb, + void* user_data) ROCPROFILER_API; + +/** + * @brief (deprecated) This call returns the number of instances specific counter contains. + * + * @param [in] agent_id rocprofiler agent identifier + * @param [in] counter_id counter id (obtained from iterate_agent_supported_counters) + * @param [out] instance_count number of instances the counter has + * @return ::rocprofiler_status_t + * @retval ROCPROFILER_STATUS_SUCCESS if counter found + * @retval ROCPROFILER_STATUS_ERROR_COUNTER_NOT_FOUND if counter not found + */ +ROCPROFILER_SDK_DEPRECATED("Information now available in rocprofiler_counter_info_v1_t. " + "This function will be removed in the future.") +rocprofiler_status_t +rocprofiler_query_counter_instance_count(rocprofiler_agent_id_t agent_id, + rocprofiler_counter_id_t counter_id, + size_t* instance_count) ROCPROFILER_API + ROCPROFILER_NONNULL(3); + +/** @} */ +ROCPROFILER_EXTERN_C_FINI diff --git a/source/include/rocprofiler-sdk/deprecated/profile_config.h b/source/include/rocprofiler-sdk/deprecated/profile_config.h new file mode 100644 index 0000000000..82c78d82cc --- /dev/null +++ b/source/include/rocprofiler-sdk/deprecated/profile_config.h @@ -0,0 +1,56 @@ +// MIT License +// +// Copyright (c) 2023-2025 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 +#include + +ROCPROFILER_EXTERN_C_INIT + +/** + * @brief (deprecated) Replaced by ::rocprofiler_create_counter_config. + * + */ +ROCPROFILER_SDK_DEPRECATED("profile_config renamed to counter_config") +static inline rocprofiler_status_t +rocprofiler_create_profile_config(rocprofiler_agent_id_t agent_id, + rocprofiler_counter_id_t* counters_list, + size_t counters_count, + rocprofiler_profile_config_id_t* config_id) +{ + return rocprofiler_create_counter_config(agent_id, counters_list, counters_count, config_id); +} + +/** + * @brief (deprecated) Replaced by ::rocprofiler_destroy_counter_config. + * + */ +ROCPROFILER_SDK_DEPRECATED("profile_config renamed to counter_config") +static inline rocprofiler_status_t +rocprofiler_destroy_profile_config(rocprofiler_profile_config_id_t config_id) +{ + return rocprofiler_destroy_counter_config(config_id); +} + +ROCPROFILER_EXTERN_C_FINI diff --git a/source/include/rocprofiler-sdk/device_counting_service.h b/source/include/rocprofiler-sdk/device_counting_service.h index 665b975da3..4d05e6440c 100644 --- a/source/include/rocprofiler-sdk/device_counting_service.h +++ b/source/include/rocprofiler-sdk/device_counting_service.h @@ -25,16 +25,17 @@ #include #include +ROCPROFILER_EXTERN_C_INIT + /** * @defgroup device_counting_service Agent Profile Counting Service * @brief needs brief description * * @{ */ -ROCPROFILER_EXTERN_C_INIT /** - * @brief Callback to set the profile config for the agent. + * @brief (experimental) Callback to set the profile config for the agent. * * @param [in] context_id context id * @param [in] config_id Profile config detailing the counters to collect for this kernel @@ -47,31 +48,33 @@ ROCPROFILER_EXTERN_C_INIT * context. * @retval ::ROCPROFILER_STATUS_SUCCESS Returned if succesfully configured */ -typedef rocprofiler_status_t (*rocprofiler_agent_set_profile_callback_t)( +ROCPROFILER_SDK_EXPERIMENTAL +typedef rocprofiler_status_t (*rocprofiler_device_counting_agent_cb_t)( rocprofiler_context_id_t context_id, - rocprofiler_profile_config_id_t config_id); + rocprofiler_counter_config_id_t config_id); /** - * @brief Configure Profile Counting Service for agent. Called when the context is started. - * Selects the counters to be used for agent profiling. + * @brief (experimental) Configure Profile Counting Service for agent. Called when the context is + * started. Selects the counters to be used for agent profiling. * * @param [in] context_id context id * @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) + * rocprofiler_device_counting_agent_cb_t) * @param [in] user_data Data supplied to rocprofiler_configure_device_counting_service */ -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); +ROCPROFILER_SDK_EXPERIMENTAL +typedef void (*rocprofiler_device_counting_service_cb_t)( + rocprofiler_context_id_t context_id, + rocprofiler_agent_id_t agent_id, + rocprofiler_device_counting_agent_cb_t set_config, + void* user_data); /** - * @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. + * @brief (experimental) 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 @@ -81,7 +84,7 @@ typedef void (*rocprofiler_device_counting_service_callback_t)( * rocprofiler_sample_device_counting_service * @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 - * counters to collect (rocprofiler_profile_config_id_t). + * counters to collect (rocprofiler_counter_config_id_t). * @param [in] user_data User supplied data to be passed to the callback cb when triggered * @return ::rocprofiler_status_t * @retval ::ROCPROFILER_STATUS_ERROR_CONTEXT_INVALID Returned if the context does not exist. @@ -90,28 +93,28 @@ typedef void (*rocprofiler_device_counting_service_callback_t)( * profiling configured for agent_id. * @retval ::ROCPROFILER_STATUS_SUCCESS Returned if succesfully configured */ +ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_status_t -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_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_cb_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_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). + * @brief (experimental) Trigger a read of the counter data for the agent profile. The counter data + * will be 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). * * @param [in] context_id context id * @param [in] user_data User supplied data, included in records outputted to buffer. * @param [in] flags Flags to specify how the counter data should be collected (defaults to sync). - * @param [in] output_records Output records collected via sampling (output is also written to - * buffer). Must be allocated by caller. - * @param [in] rec_count On entry, this is the maximum number of records rocprof can store in - * output_records. On exit, contains the number of actual records. + * @param [in] output_records (Optional) Provides the values immediately instead of outputting to + * buffer. Must be allocated by caller. + * @param [in] rec_count (Optional) On entry, this is the maximum number of records rocprof can + * store in output_records. On exit, contains the number of actual records. * @return ::rocprofiler_status_t * @retval ::ROCPROFILER_STATUS_ERROR_CONTEXT_INVALID Returned if the context does not exist or * the context is not configured for agent profiling. @@ -124,6 +127,7 @@ rocprofiler_configure_device_counting_service(rocprofiler_context_id_t context_i * @retval ::ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT Returned If ASYNC is being used while * output_records is not null. */ +ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_status_t rocprofiler_sample_device_counting_service(rocprofiler_context_id_t context_id, rocprofiler_user_data_t user_data, diff --git a/source/include/rocprofiler-sdk/dispatch_counting_service.h b/source/include/rocprofiler-sdk/dispatch_counting_service.h index cd5c4cc2ea..dfb1fe8c31 100644 --- a/source/include/rocprofiler-sdk/dispatch_counting_service.h +++ b/source/include/rocprofiler-sdk/dispatch_counting_service.h @@ -23,10 +23,10 @@ #pragma once #include +#include #include #include #include -#include ROCPROFILER_EXTERN_C_INIT @@ -38,10 +38,10 @@ ROCPROFILER_EXTERN_C_INIT */ /** - * @brief Kernel dispatch data for profile counting callbacks. + * @brief (experimental) Kernel dispatch data for profile counting callbacks. * */ -typedef struct rocprofiler_dispatch_counting_service_data_t +typedef struct ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_dispatch_counting_service_data_t { uint64_t size; ///< Size of this struct rocprofiler_async_correlation_id_t correlation_id; ///< Correlation ID for this dispatch @@ -51,11 +51,11 @@ typedef struct rocprofiler_dispatch_counting_service_data_t } rocprofiler_dispatch_counting_service_data_t; /** - * @brief ROCProfiler Profile Counting Counter Record Header Information + * @brief (experimental) ROCProfiler Profile Counting Counter Record Header Information * * This is buffer equivalent of ::rocprofiler_dispatch_counting_service_data_t */ -typedef struct rocprofiler_dispatch_counting_service_record_t +typedef struct ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_dispatch_counting_service_record_t { uint64_t size; ///< Size of this struct uint64_t num_records; ///< number of ::rocprofiler_record_counter_t records @@ -66,27 +66,27 @@ typedef struct rocprofiler_dispatch_counting_service_record_t } rocprofiler_dispatch_counting_service_record_t; /** - * @brief Kernel Dispatch Callback. This is a callback that is invoked before the kernel - * is enqueued into the HSA queue. What counters to collect for a kernel are set - * via passing back a profile config (config) in this callback. These counters - * will be collected and emplaced in the buffer with @ref rocprofiler_buffer_id_t used when - * setting up this callback. + * @brief (experimental) Kernel Dispatch Callback. This is a callback that is invoked before the + * kernel is enqueued into the HSA queue. What counters to collect for a kernel are set via passing + * back a profile config (config) in this callback. These counters 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_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_counting_service */ -typedef void (*rocprofiler_dispatch_counting_service_callback_t)( +ROCPROFILER_SDK_EXPERIMENTAL +typedef void (*rocprofiler_dispatch_counting_service_cb_t)( rocprofiler_dispatch_counting_service_data_t dispatch_data, - rocprofiler_profile_config_id_t* config, + rocprofiler_counter_config_id_t* config, rocprofiler_user_data_t* user_data, void* callback_data_args); /** - * @brief Counting record callback. This is a callback is invoked when the kernel + * @brief (experimental) 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_dispatch_counting_service_callback_t. Only used with + * @ref rocprofiler_dispatch_counting_service_cb_t. Only used with * @ref rocprofiler_configure_callback_dispatch_counting_service. * * @param [in] dispatch_data @see ::rocprofiler_dispatch_counting_service_data_t @@ -95,7 +95,8 @@ typedef void (*rocprofiler_dispatch_counting_service_callback_t)( * @param [in] user_data User data instance from dispatch callback * @param [in] callback_data_args Callback supplied via buffered_dispatch_counting_service */ -typedef void (*rocprofiler_profile_counting_record_callback_t)( +ROCPROFILER_SDK_EXPERIMENTAL +typedef void (*rocprofiler_dispatch_counting_record_cb_t)( rocprofiler_dispatch_counting_service_data_t dispatch_data, rocprofiler_record_counter_t* record_data, size_t record_count, @@ -103,7 +104,7 @@ typedef void (*rocprofiler_profile_counting_record_callback_t)( void* callback_data_args); /** - * @brief Configure buffered dispatch profile Counting Service. + * @brief (experimental) Configure buffered dispatch profile Counting Service. * Collects the counters in dispatch packets and stores them * in a buffer with @p buffer_id. The buffer may contain packets from more than * one dispatch (denoted by correlation id). Will trigger the @@ -112,14 +113,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_dispatch_counting_service_callback_t (i.e. + * @ref rocprofiler_dispatch_counting_service_cb_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_dispatch_counting_service_callback_t (i.e. + * @ref rocprofiler_dispatch_counting_service_cb_t (i.e. * tool code checking the agent param to see if they want to profile * it). * @@ -129,15 +130,16 @@ typedef void (*rocprofiler_profile_counting_record_callback_t)( * @param [in] callback_data_args callback data * @return ::rocprofiler_status_t */ +ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_status_t -rocprofiler_configure_buffered_dispatch_counting_service( - rocprofiler_context_id_t context_id, - rocprofiler_buffer_id_t buffer_id, - rocprofiler_dispatch_counting_service_callback_t callback, - void* callback_data_args) ROCPROFILER_API; +rocprofiler_configure_buffer_dispatch_counting_service( + rocprofiler_context_id_t context_id, + rocprofiler_buffer_id_t buffer_id, + rocprofiler_dispatch_counting_service_cb_t callback, + void* callback_data_args) ROCPROFILER_API; /** - * @brief Configure buffered dispatch profile Counting Service. + * @brief (experimental) Configure buffered dispatch profile Counting Service. * Collects the counters in dispatch packets and calls a callback * with the counters collected during that dispatch. * @@ -148,13 +150,15 @@ rocprofiler_configure_buffered_dispatch_counting_service( * @param [in] record_callback_args Callback args for record callback * @return ::rocprofiler_status_t */ +ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_status_t rocprofiler_configure_callback_dispatch_counting_service( - rocprofiler_context_id_t context_id, - 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; + rocprofiler_context_id_t context_id, + rocprofiler_dispatch_counting_service_cb_t dispatch_callback, + void* dispatch_callback_args, + rocprofiler_dispatch_counting_record_cb_t record_callback, + void* record_callback_args) ROCPROFILER_API; + /** @} */ ROCPROFILER_EXTERN_C_FINI diff --git a/source/include/rocprofiler-sdk/experimental/counters.h b/source/include/rocprofiler-sdk/experimental/counters.h index bd5c460af2..56fea6ed0a 100644 --- a/source/include/rocprofiler-sdk/experimental/counters.h +++ b/source/include/rocprofiler-sdk/experimental/counters.h @@ -29,10 +29,10 @@ ROCPROFILER_EXTERN_C_INIT /** - * @brief Apply a custom counter definition (YAML). This function must be called before - * counter iteration functions like @ref rocprofiler_iterate_agent_supported_counters - * if custom counter definitions are to be used. This function will return an error - * if the definition has already been loaded. + * @brief (experimental) Apply a custom counter definition (YAML). This function must be called + * before counter iteration functions like @ref rocprofiler_iterate_agent_supported_counters if + * custom counter definitions are to be used. This function will return an error if the definition + * has already been loaded. * @param [in] yaml YAML string containing counter definitions * @param [in] size Size of the YAML string * @param [in] flags Flags to apply to the counter definition @@ -40,6 +40,7 @@ ROCPROFILER_EXTERN_C_INIT * @retval ROCPROFILER_STATUS_SUCCESS if counter definition applied * @retval ROCPROFILER_STATUS_ERROR if counter definition already loaded */ +ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_status_t rocprofiler_load_counter_definition(const char* yaml, size_t size, diff --git a/source/include/rocprofiler-sdk/ext_version.h.in b/source/include/rocprofiler-sdk/ext_version.h.in new file mode 100644 index 0000000000..99b9186062 --- /dev/null +++ b/source/include/rocprofiler-sdk/ext_version.h.in @@ -0,0 +1,55 @@ +// MIT License +// +// Copyright (c) 2023-2025 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 + +/** + * @file rocprofiler-sdk/ext_version.h + * @brief Defines versions for external dependencies + * + */ + +// clang-format off +#define ROCPROFILER_HSA_RUNTIME_VERSION_MAJOR @HSA_RUNTIME_VERSION_MAJOR@ +#define ROCPROFILER_HSA_RUNTIME_VERSION_MINOR @HSA_RUNTIME_VERSION_MINOR@ + +#cmakedefine ROCPROFILER_HSA_API_TABLE_MAJOR_VERSION @ROCPROFILER_HSA_API_TABLE_MAJOR_VERSION@ +#cmakedefine ROCPROFILER_HSA_CORE_API_TABLE_MAJOR_VERSION @ROCPROFILER_HSA_CORE_API_TABLE_MAJOR_VERSION@ +#cmakedefine ROCPROFILER_HSA_AMD_EXT_API_TABLE_MAJOR_VERSION @ROCPROFILER_HSA_AMD_EXT_API_TABLE_MAJOR_VERSION@ +#cmakedefine ROCPROFILER_HSA_FINALIZER_API_TABLE_MAJOR_VERSION @ROCPROFILER_HSA_FINALIZER_API_TABLE_MAJOR_VERSION@ +#cmakedefine ROCPROFILER_HSA_IMAGE_API_TABLE_MAJOR_VERSION @ROCPROFILER_HSA_IMAGE_API_TABLE_MAJOR_VERSION@ +#cmakedefine ROCPROFILER_HSA_AQLPROFILE_API_TABLE_MAJOR_VERSION @ROCPROFILER_HSA_AQLPROFILE_API_TABLE_MAJOR_VERSION@ +#cmakedefine ROCPROFILER_HSA_TOOLS_API_TABLE_MAJOR_VERSION @ROCPROFILER_HSA_TOOLS_API_TABLE_MAJOR_VERSION@ + +#cmakedefine ROCPROFILER_HSA_API_TABLE_STEP_VERSION @ROCPROFILER_HSA_API_TABLE_STEP_VERSION@ +#cmakedefine ROCPROFILER_HSA_CORE_API_TABLE_STEP_VERSION @ROCPROFILER_HSA_CORE_API_TABLE_STEP_VERSION@ +#cmakedefine ROCPROFILER_HSA_AMD_EXT_API_TABLE_STEP_VERSION @ROCPROFILER_HSA_AMD_EXT_API_TABLE_STEP_VERSION@ +#cmakedefine ROCPROFILER_HSA_FINALIZER_API_TABLE_STEP_VERSION @ROCPROFILER_HSA_FINALIZER_API_TABLE_STEP_VERSION@ +#cmakedefine ROCPROFILER_HSA_IMAGE_API_TABLE_STEP_VERSION @ROCPROFILER_HSA_IMAGE_API_TABLE_STEP_VERSION@ +#cmakedefine ROCPROFILER_HSA_AQLPROFILE_API_TABLE_STEP_VERSION @ROCPROFILER_HSA_AQLPROFILE_API_TABLE_STEP_VERSION@ +#cmakedefine ROCPROFILER_HSA_TOOLS_API_TABLE_STEP_VERSION @ROCPROFILER_HSA_TOOLS_API_TABLE_STEP_VERSION@ +// clang-format on + +// latest hsa-runtime version supported +#define ROCPROFILER_HSA_RUNTIME_VERSION \ + ((10000 * ROCPROFILER_HSA_RUNTIME_VERSION_MAJOR) + \ + (100 * ROCPROFILER_HSA_RUNTIME_VERSION_MINOR)) diff --git a/source/include/rocprofiler-sdk/external_correlation.h b/source/include/rocprofiler-sdk/external_correlation.h index 54898843f5..9a6ae17af4 100644 --- a/source/include/rocprofiler-sdk/external_correlation.h +++ b/source/include/rocprofiler-sdk/external_correlation.h @@ -35,7 +35,7 @@ ROCPROFILER_EXTERN_C_INIT */ /** - * @brief ROCProfiler External Correlation ID Operations. + * @brief (experimental) ROCProfiler External Correlation ID Operations. * * These kinds correspond to callback and buffered tracing kinds (@see * ::rocprofiler_callback_tracing_kind_t and ::rocprofiler_buffer_tracing_kind_t) which generate @@ -51,7 +51,8 @@ ROCPROFILER_EXTERN_C_INIT * where the external correlation ID value is not important while also getting a request for an * external correlation ID for other tracing kinds. */ -typedef enum // NOLINT(performance-enum-size) +// NOLINTNEXTLINE(performance-enum-size) +typedef enum ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_external_correlation_id_request_kind_t { ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_NONE = 0, ///< Unknown kind ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_HSA_CORE_API, ///< @@ -75,7 +76,7 @@ typedef enum // NOLINT(performance-enum-size) } rocprofiler_external_correlation_id_request_kind_t; /** - * @brief Callback requesting a value for the external correlation id. + * @brief (experimental) Callback requesting a value for the external correlation id. * * @param [in] thread_id Id of the thread making the request * @param [in] context_id Id of the context making the request @@ -95,6 +96,7 @@ typedef enum // NOLINT(performance-enum-size) * correlation ID value and the thread-local value for the most recently pushed external correlation * ID should be used instead */ +ROCPROFILER_SDK_EXPERIMENTAL typedef int (*rocprofiler_external_correlation_id_request_cb_t)( rocprofiler_thread_id_t thread_id, rocprofiler_context_id_t context_id, @@ -105,7 +107,7 @@ typedef int (*rocprofiler_external_correlation_id_request_cb_t)( void* data); /** - * @brief Configure External Correlation ID Request Service. + * @brief (experimental) Configure External Correlation ID Request Service. * * @param [in] context_id Context to associate the service with * @param [in] kinds Array of ::rocprofiler_external_correlation_id_request_kind_t values. If @@ -117,13 +119,14 @@ typedef int (*rocprofiler_external_correlation_id_request_cb_t)( * @param [in] callback_args Data provided to every invocation of the callback function * @return ::rocprofiler_status_t * @retval ::ROCPROFILER_STATUS_ERROR_CONFIGURATION_LOCKED Invoked outside of the initialization - * function in @ref rocprofiler_tool_configure_result_t provided to rocprofiler via @ref - * rocprofiler_configure function + * function in ::rocprofiler_tool_configure_result_t provided to rocprofiler via + * ::rocprofiler_configure function * @retval ::ROCPROFILER_STATUS_ERROR_CONTEXT_NOT_FOUND The provided context is not valid/registered - * @retval ::ROCPROFILER_STATUS_ERROR_SERVICE_ALREADY_CONFIGURED if the same @ref - * rocprofiler_callback_tracing_kind_t value is provided more than once (per context) -- in - * other words, we do not support overriding or combining the kinds in separate function calls. + * @retval ::ROCPROFILER_STATUS_ERROR_SERVICE_ALREADY_CONFIGURED if the same + * ::rocprofiler_callback_tracing_kind_t value is provided more than once (per context) -- in other + * words, we do not support overriding or combining the kinds in separate function calls. */ +ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_status_t rocprofiler_configure_external_correlation_id_request_service( rocprofiler_context_id_t context_id, @@ -133,7 +136,7 @@ rocprofiler_configure_external_correlation_id_request_service( void* callback_args) ROCPROFILER_API ROCPROFILER_NONNULL(4); /** - * @brief Push default value for `external` field in @ref rocprofiler_correlation_id_t onto stack. + * @brief Push default value for `external` field in ::rocprofiler_correlation_id_t onto stack. * * External correlation ids are thread-local values. However, if rocprofiler internally requests an * external correlation id on a non-main thread and an external correlation id has not been pushed @@ -144,8 +147,8 @@ rocprofiler_configure_external_correlation_id_request_service( * * @param [in] context Associated context * @param [in] tid thread identifier. @see rocprofiler_get_thread_id - * @param [in] external_correlation_id User data to place in external field in @ref - * rocprofiler_correlation_id_t + * @param [in] external_correlation_id User data to place in external field in + * ::rocprofiler_correlation_id_t * @return ::rocprofiler_status_t * @retval ::ROCPROFILER_STATUS_ERROR_CONTEXT_NOT_FOUND Context does not exist * @retval ::ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT Thread id is not valid @@ -157,7 +160,7 @@ rocprofiler_push_external_correlation_id(rocprofiler_context_id_t context, ROCPROFILER_API; /** - * @brief Pop default value for `external` field in @ref rocprofiler_correlation_id_t off of stack + * @brief Pop default value for `external` field in ::rocprofiler_correlation_id_t off of stack. * * @param [in] context Associated context * @param [in] tid thread identifier. @see rocprofiler_get_thread_id diff --git a/source/include/rocprofiler-sdk/fwd.h b/source/include/rocprofiler-sdk/fwd.h index e60e411974..a97eea9dba 100644 --- a/source/include/rocprofiler-sdk/fwd.h +++ b/source/include/rocprofiler-sdk/fwd.h @@ -42,14 +42,10 @@ ROCPROFILER_EXTERN_C_INIT * @{ */ -// TODO(aelwazir): Do we need to add a null (way) for every handle? -// TODO(aelwazir): Remove API Data args from the doxygen? -// TODO(aelwazir): Not everything in bin needs to be installed bin, use libexec or share? - /** * @brief Status codes. */ -typedef enum // NOLINT(performance-enum-size) +typedef enum rocprofiler_status_t // NOLINT(performance-enum-size) { ROCPROFILER_STATUS_SUCCESS = 0, ///< No error occurred ROCPROFILER_STATUS_ERROR, ///< Generalized error @@ -111,10 +107,10 @@ typedef enum // NOLINT(performance-enum-size) } rocprofiler_status_t; /** - * @brief Buffer record categories. This enumeration type is encoded in @ref - * rocprofiler_record_header_t category field + * @brief Buffer record categories. This enumeration type is encoded in + * ::rocprofiler_record_header_t category field */ -typedef enum // NOLINT(performance-enum-size) +typedef enum rocprofiler_buffer_category_t // NOLINT(performance-enum-size) { ROCPROFILER_BUFFER_CATEGORY_NONE = 0, ROCPROFILER_BUFFER_CATEGORY_TRACING, @@ -126,7 +122,7 @@ typedef enum // NOLINT(performance-enum-size) /** * @brief Agent type. */ -typedef enum // NOLINT(performance-enum-size) +typedef enum rocprofiler_agent_type_t // NOLINT(performance-enum-size) { ROCPROFILER_AGENT_TYPE_NONE = 0, ///< Agent type is unknown ROCPROFILER_AGENT_TYPE_CPU, ///< Agent type is a CPU @@ -137,7 +133,7 @@ typedef enum // NOLINT(performance-enum-size) /** * @brief Service Callback Phase. */ -typedef enum // NOLINT(performance-enum-size) +typedef enum rocprofiler_callback_phase_t // NOLINT(performance-enum-size) { ROCPROFILER_CALLBACK_PHASE_NONE = 0, ///< Callback has no phase ROCPROFILER_CALLBACK_PHASE_ENTER, ///< Callback invoked prior to function execution @@ -152,7 +148,7 @@ typedef enum // NOLINT(performance-enum-size) /** * @brief Service Callback Tracing Kind. @see rocprofiler_configure_callback_tracing_service. */ -typedef enum // NOLINT(performance-enum-size) +typedef enum rocprofiler_callback_tracing_kind_t // NOLINT(performance-enum-size) { ROCPROFILER_CALLBACK_TRACING_NONE = 0, ROCPROFILER_CALLBACK_TRACING_HSA_CORE_API, ///< @see ::rocprofiler_hsa_core_api_id_t @@ -178,15 +174,14 @@ typedef enum // NOLINT(performance-enum-size) ///< library has been initialized ROCPROFILER_CALLBACK_TRACING_ROCDECODE_API, ///< rocDecode API Tracing ROCPROFILER_CALLBACK_TRACING_ROCJPEG_API, ///< rocJPEG API Tracing - ROCPROFILER_CALLBACK_TRACING_HIP_STREAM_API, ///< @see - ///< ::rocprofiler_hip_stream_operation_t + ROCPROFILER_CALLBACK_TRACING_HIP_STREAM, ///< @see ::rocprofiler_hip_stream_operation_t ROCPROFILER_CALLBACK_TRACING_LAST, } rocprofiler_callback_tracing_kind_t; /** * @brief Service Buffer Tracing Kind. @see rocprofiler_configure_buffer_tracing_service. */ -typedef enum // NOLINT(performance-enum-size) +typedef enum rocprofiler_buffer_tracing_kind_t // NOLINT(performance-enum-size) { ROCPROFILER_BUFFER_TRACING_NONE = 0, ROCPROFILER_BUFFER_TRACING_HSA_CORE_API, ///< @see ::rocprofiler_hsa_core_api_id_t @@ -211,9 +206,9 @@ typedef enum // NOLINT(performance-enum-size) ROCPROFILER_BUFFER_TRACING_RUNTIME_INITIALIZATION, ///< Record indicating a runtime library has ///< been initialized. @see ///< ::rocprofiler_runtime_initialization_operation_t - ROCPROFILER_BUFFER_TRACING_ROCDECODE_API, ///< rocDecode tracing - ROCPROFILER_BUFFER_TRACING_ROCJPEG_API, ///< rocJPEG tracing - ROCPROFILER_BUFFER_TRACING_HIP_STREAM_API, ///< Display HIP Stream + ROCPROFILER_BUFFER_TRACING_ROCDECODE_API, ///< rocDecode tracing + ROCPROFILER_BUFFER_TRACING_ROCJPEG_API, ///< rocJPEG tracing + ROCPROFILER_BUFFER_TRACING_HIP_STREAM, ///< @see ::rocprofiler_hip_stream_operation_t ROCPROFILER_BUFFER_TRACING_HIP_RUNTIME_API_EXT, ROCPROFILER_BUFFER_TRACING_HIP_COMPILER_API_EXT, ROCPROFILER_BUFFER_TRACING_LAST, @@ -229,7 +224,7 @@ typedef enum // NOLINT(performance-enum-size) /** * @brief ROCProfiler Code Object Tracer Operations. */ -typedef enum // NOLINT(performance-enum-size) +typedef enum rocprofiler_code_object_operation_t // NOLINT(performance-enum-size) { ROCPROFILER_CODE_OBJECT_NONE = 0, ///< Unknown code object operation ROCPROFILER_CODE_OBJECT_LOAD, ///< Code object containing kernel symbols @@ -239,9 +234,10 @@ typedef enum // NOLINT(performance-enum-size) } rocprofiler_code_object_operation_t; /** - * @brief ROCProfiler Stream Handle Operations. + * @brief ROCProfiler HIP Stream Operations. These operations can be used to associate subsequent + * information with a HIP stream */ -typedef enum // NOLINT(performance-enum-size) +typedef enum rocprofiler_hip_stream_operation_t // NOLINT(performance-enum-size) { ROCPROFILER_HIP_STREAM_NONE = 0, ///< Unknown stream handle operation ROCPROFILER_HIP_STREAM_CREATE, ///< A stream handle is created @@ -258,7 +254,7 @@ typedef enum // NOLINT(performance-enum-size) /** * @brief Memory Copy Operations. */ -typedef enum // NOLINT(performance-enum-size) +typedef enum rocprofiler_memory_copy_operation_t // NOLINT(performance-enum-size) { ROCPROFILER_MEMORY_COPY_NONE = 0, ///< Unknown memory copy direction ROCPROFILER_MEMORY_COPY_HOST_TO_HOST, ///< Memory copy from host to host @@ -271,7 +267,7 @@ typedef enum // NOLINT(performance-enum-size) /** * @brief Memory Allocation Operation. */ -typedef enum // NOLINT(performance-enum-size) +typedef enum rocprofiler_memory_allocation_operation_t // NOLINT(performance-enum-size) { ROCPROFILER_MEMORY_ALLOCATION_NONE = 0, ///< Unknown memory allocation function ROCPROFILER_MEMORY_ALLOCATION_ALLOCATE, ///< Allocate memory function @@ -284,7 +280,7 @@ typedef enum // NOLINT(performance-enum-size) /** * @brief ROCProfiler Kernel Dispatch Tracing Operation Types. */ -typedef enum // NOLINT(performance-enum-size) +typedef enum rocprofiler_kernel_dispatch_operation_t // NOLINT(performance-enum-size) { ROCPROFILER_KERNEL_DISPATCH_NONE = 0, ///< Unknown kernel dispatch operation ROCPROFILER_KERNEL_DISPATCH_ENQUEUE = 1, @@ -330,7 +326,7 @@ typedef enum // NOLINT(performance-enum-size) /** * @brief PC Sampling Method. */ -typedef enum // NOLINT(performance-enum-size) +typedef enum rocprofiler_pc_sampling_method_t // NOLINT(performance-enum-size) { ROCPROFILER_PC_SAMPLING_METHOD_NONE = 0, ///< Unknown sampling type ROCPROFILER_PC_SAMPLING_METHOD_STOCHASTIC, ///< Stochastic sampling (MI300+) @@ -341,7 +337,7 @@ typedef enum // NOLINT(performance-enum-size) /** * @brief PC Sampling Unit. */ -typedef enum // NOLINT(performance-enum-size) +typedef enum rocprofiler_pc_sampling_unit_t // NOLINT(performance-enum-size) { ROCPROFILER_PC_SAMPLING_UNIT_NONE = 0, ///< Sample interval has unspecified units ROCPROFILER_PC_SAMPLING_UNIT_INSTRUCTIONS, ///< Sample interval is in instructions @@ -353,7 +349,7 @@ typedef enum // NOLINT(performance-enum-size) /** * @brief Actions when Buffer is full. */ -typedef enum // NOLINT(performance-enum-size) +typedef enum rocprofiler_buffer_policy_t // NOLINT(performance-enum-size) { ROCPROFILER_BUFFER_POLICY_NONE = 0, ///< No policy has been set ROCPROFILER_BUFFER_POLICY_DISCARD, ///< Drop records when buffer is full @@ -364,7 +360,7 @@ typedef enum // NOLINT(performance-enum-size) /** * @brief Page migration event. */ -typedef enum // NOLINT(performance-enum-size) +typedef enum rocprofiler_page_migration_operation_t // NOLINT(performance-enum-size) { ROCPROFILER_PAGE_MIGRATION_NONE = 0, ///< Unknown event ROCPROFILER_PAGE_MIGRATION_PAGE_MIGRATE_START, @@ -381,7 +377,7 @@ typedef enum // NOLINT(performance-enum-size) /** * @brief Scratch event kind */ -typedef enum +typedef enum rocprofiler_scratch_memory_operation_t { ROCPROFILER_SCRATCH_MEMORY_NONE = 0, ///< Unknown scratch operation ROCPROFILER_SCRATCH_MEMORY_ALLOC, ///< Scratch memory allocation event @@ -394,7 +390,7 @@ typedef enum * @brief Enumeration for specifying runtime libraries supported by rocprofiler. This enumeration is * used for thread creation callbacks. @see INTERNAL_THREADING. */ -typedef enum +typedef enum rocprofiler_runtime_library_t { ROCPROFILER_LIBRARY = (1 << 0), ROCPROFILER_HSA_LIBRARY = (1 << 1), @@ -410,7 +406,7 @@ typedef enum * @brief Enumeration for specifying intercept tables supported by rocprofiler. This enumeration is * used for intercept tables. @see INTERCEPT_TABLE. */ -typedef enum +typedef enum rocprofiler_intercept_table_t { ROCPROFILER_HSA_TABLE = (1 << 0), ROCPROFILER_HIP_RUNTIME_TABLE = (1 << 1), @@ -427,7 +423,7 @@ typedef enum /** * @brief ROCProfiler Runtime Initialization Tracer Operations. */ -typedef enum // NOLINT(performance-enum-size) +typedef enum rocprofiler_runtime_initialization_operation_t // NOLINT(performance-enum-size) { ROCPROFILER_RUNTIME_INITIALIZATION_NONE = 0, ///< Unknown runtime initialization ROCPROFILER_RUNTIME_INITIALIZATION_HSA, ///< Application loaded HSA runtime @@ -442,10 +438,11 @@ typedef enum // NOLINT(performance-enum-size) /** * @brief Enumeration for specifying the counter info struct version you want. */ -typedef enum +typedef enum rocprofiler_counter_info_version_id_t { ROCPROFILER_COUNTER_INFO_VERSION_NONE, ROCPROFILER_COUNTER_INFO_VERSION_0, ///< @see ::rocprofiler_counter_info_v0_t + ROCPROFILER_COUNTER_INFO_VERSION_1, ///< @see ::rocprofiler_counter_info_v1_t ROCPROFILER_COUNTER_INFO_VERSION_LAST, } rocprofiler_counter_info_version_id_t; @@ -453,7 +450,7 @@ typedef enum * @brief Enumeration for distinguishing different buffer record kinds within the * ::ROCPROFILER_BUFFER_CATEGORY_COUNTERS category */ -typedef enum +typedef enum rocprofiler_counter_record_kind_t { ROCPROFILER_COUNTER_RECORD_NONE = 0, ROCPROFILER_COUNTER_RECORD_PROFILE_COUNTING_DISPATCH_HEADER, ///< ::rocprofiler_dispatch_counting_service_record_t @@ -468,7 +465,7 @@ typedef enum /** * @brief Enumeration of flags that can be used with some counter api calls */ -typedef enum +typedef enum rocprofiler_counter_flag_t { ROCPROFILER_COUNTER_FLAG_NONE = 0, ROCPROFILER_COUNTER_FLAG_ASYNC, ///< Do not wait for completion before returning. @@ -481,7 +478,7 @@ typedef enum * @brief Enumeration for distinguishing different buffer record kinds within the * ::ROCPROFILER_BUFFER_CATEGORY_PC_SAMPLING category */ -typedef enum +typedef enum rocprofiler_pc_sampling_record_kind_t { ROCPROFILER_PC_SAMPLING_RECORD_NONE = 0, ROCPROFILER_PC_SAMPLING_RECORD_INVALID_SAMPLE, ///< ::rocprofiler_pc_sampling_record_invalid_t @@ -584,6 +581,16 @@ typedef union rocprofiler_uuid_t // //--------------------------------------------------------------------------------------// +/** + * @brief Versioning info. + */ +typedef struct rocprofiler_version_triplet_t +{ + uint32_t major; + uint32_t minor; + uint32_t patch; +} rocprofiler_version_triplet_t; + /** * @brief Context ID. */ @@ -679,12 +686,12 @@ typedef struct rocprofiler_counter_id_t /** * @brief Profile Configurations - * @see rocprofiler_create_profile_config for how to create. + * @see rocprofiler_create_counter_config for how to create. */ -typedef struct rocprofiler_profile_config_id_t +typedef struct rocprofiler_counter_config_id_t { uint64_t handle; // Opaque handle -} rocprofiler_profile_config_id_t; +} rocprofiler_counter_config_id_t; /** * @brief Multi-dimensional struct of data used to describe GPU workgroup and grid sizes @@ -740,13 +747,13 @@ typedef struct rocprofiler_callback_tracing_record_t * * @endcode */ -typedef struct +typedef struct rocprofiler_record_header_t { union { struct { - uint32_t category; ///< rocprofiler_buffer_category_t + uint32_t category; ///< ::rocprofiler_buffer_category_t uint32_t kind; ///< domain }; uint64_t hash; ///< generic identifier. You can compute this via: `uint64_t hash = category @@ -756,12 +763,12 @@ typedef struct } rocprofiler_record_header_t; /** - * @brief Function for computing the unsigned 64-bit hash value in @ref rocprofiler_record_header_t + * @brief Function for computing the unsigned 64-bit hash value in ::rocprofiler_record_header_t * from a category and kind (two unsigned 32-bit values) * - * @param [in] category a value from @ref rocprofiler_buffer_category_t - * @param [in] kind depending on the category, this is the domain value, e.g., @ref - * rocprofiler_buffer_tracing_kind_t value + * @param [in] category a value from ::rocprofiler_buffer_category_t + * @param [in] kind depending on the category, this is the domain value, e.g., + * ::rocprofiler_buffer_tracing_kind_t value * @return uint64_t hash value of category and kind */ static inline uint64_t @@ -801,20 +808,22 @@ typedef struct rocprofiler_kernel_dispatch_info_t /** * @brief Details for the dimension, including its size, for a counter record. */ -typedef struct +typedef struct rocprofiler_counter_record_dimension_info_t { const char* name; size_t instance_size; rocprofiler_counter_dimension_id_t id; /// @var id - /// @brief Id for this dimension used by @ref rocprofiler_query_record_dimension_position -} rocprofiler_record_dimension_info_t; + /// @brief Id for this dimension used by ::rocprofiler_query_record_dimension_position +} rocprofiler_counter_record_dimension_info_t; + +typedef rocprofiler_counter_record_dimension_info_t rocprofiler_record_dimension_info_t; /** * @brief ROCProfiler Profile Counting Counter Record per instance. */ -typedef struct +typedef struct rocprofiler_counter_record_t { rocprofiler_counter_instance_id_t id; ///< counter identifier double counter_value; ///< counter value @@ -831,27 +840,15 @@ typedef struct /// instance (provided during callback for profile config) or a /// ::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; +} rocprofiler_counter_record_t; -/** - * @brief Counter info struct version 0 - */ -typedef struct -{ - rocprofiler_counter_id_t id; ///< Id of this counter - const char* name; ///< Name of the counter - const char* description; ///< Description of the counter - const char* block; ///< Block of the counter (non-derived only) - const char* expression; ///< Counter expression (derived counters only) - uint8_t is_constant : 1; ///< If this counter is HW constant - uint8_t is_derived : 1; ///< If this counter is a derived counter -} rocprofiler_counter_info_v0_t; +typedef rocprofiler_counter_record_t rocprofiler_record_counter_t; /** * @brief ROCProfiler SPM Record. * */ -typedef struct +typedef struct rocprofiler_spm_record_t { /** * Counters, including identifiers to get counter information and Counters @@ -861,6 +858,14 @@ typedef struct uint64_t counters_count; } rocprofiler_spm_record_t; +#if defined(ROCPROFILER_SDK_BETA_COMPAT) && ROCPROFILER_SDK_BETA_COMPAT > 0 + +// "profile_config" renamed to "counter_config" +ROCPROFILER_SDK_DEPRECATED("profile_config renamed to counter_config") +typedef rocprofiler_counter_config_id_t rocprofiler_profile_config_id_t; + +#endif + /** @} */ ROCPROFILER_EXTERN_C_FINI diff --git a/source/include/rocprofiler-sdk/hip/api_args.h b/source/include/rocprofiler-sdk/hip/api_args.h index 0a3657d82d..b86105c77d 100644 --- a/source/include/rocprofiler-sdk/hip/api_args.h +++ b/source/include/rocprofiler-sdk/hip/api_args.h @@ -24,7 +24,6 @@ #include #include -#include #include #include diff --git a/source/include/rocprofiler-sdk/hip/compiler_api_id.h b/source/include/rocprofiler-sdk/hip/compiler_api_id.h index f72408e85d..e44c1533cc 100644 --- a/source/include/rocprofiler-sdk/hip/compiler_api_id.h +++ b/source/include/rocprofiler-sdk/hip/compiler_api_id.h @@ -22,14 +22,12 @@ #pragma once -#include - #include /** * @brief ROCProfiler enumeration of HIP Compiler API tracing operations */ -typedef enum // NOLINT(performance-enum-size) +typedef enum rocprofiler_hip_compiler_api_id_t // NOLINT(performance-enum-size) { ROCPROFILER_HIP_COMPILER_API_ID_NONE = -1, ROCPROFILER_HIP_COMPILER_API_ID___hipPopCallConfiguration = 0, diff --git a/source/include/rocprofiler-sdk/hip/runtime_api_id.h b/source/include/rocprofiler-sdk/hip/runtime_api_id.h index c7f6d262ab..16d2cdc910 100644 --- a/source/include/rocprofiler-sdk/hip/runtime_api_id.h +++ b/source/include/rocprofiler-sdk/hip/runtime_api_id.h @@ -22,14 +22,12 @@ #pragma once -#include - #include /** * @brief ROCProfiler enumeration of HIP runtime API tracing operations */ -typedef enum // NOLINT(performance-enum-size) +typedef enum rocprofiler_hip_runtime_api_id_t // NOLINT(performance-enum-size) { ROCPROFILER_HIP_RUNTIME_API_ID_NONE = -1, ROCPROFILER_HIP_RUNTIME_API_ID_hipApiName = 0, diff --git a/source/include/rocprofiler-sdk/hip/table_id.h b/source/include/rocprofiler-sdk/hip/table_id.h index 45cd2c6212..1051eb666f 100644 --- a/source/include/rocprofiler-sdk/hip/table_id.h +++ b/source/include/rocprofiler-sdk/hip/table_id.h @@ -23,7 +23,7 @@ #pragma once // NOLINTNEXTLINE(performance-enum-size) -typedef enum +typedef enum rocprofiler_hip_table_id_t { ROCPROFILER_HIP_TABLE_ID_NONE = -1, ROCPROFILER_HIP_TABLE_ID_Compiler = 0, diff --git a/source/include/rocprofiler-sdk/hsa/amd_ext_api_id.h b/source/include/rocprofiler-sdk/hsa/amd_ext_api_id.h index 9e05765596..ae95d45599 100644 --- a/source/include/rocprofiler-sdk/hsa/amd_ext_api_id.h +++ b/source/include/rocprofiler-sdk/hsa/amd_ext_api_id.h @@ -22,13 +22,13 @@ #pragma once +#include #include -#include /** * @brief ROCProfiler enumeration of HSA AMD Extended API tracing operations */ -typedef enum // NOLINT(performance-enum-size) +typedef enum rocprofiler_hsa_amd_ext_api_id_t // NOLINT(performance-enum-size) { ROCPROFILER_HSA_AMD_EXT_API_ID_NONE = -1, ROCPROFILER_HSA_AMD_EXT_API_ID_hsa_amd_coherency_get_type, diff --git a/source/include/rocprofiler-sdk/hsa/api_args.h b/source/include/rocprofiler-sdk/hsa/api_args.h index 704d443715..bfc70c8c0d 100644 --- a/source/include/rocprofiler-sdk/hsa/api_args.h +++ b/source/include/rocprofiler-sdk/hsa/api_args.h @@ -23,8 +23,8 @@ #pragma once #include +#include #include -#include #include #include diff --git a/source/include/rocprofiler-sdk/hsa/api_trace_version.h b/source/include/rocprofiler-sdk/hsa/api_trace_version.h index 7241de8b55..39dec1404a 100644 --- a/source/include/rocprofiler-sdk/hsa/api_trace_version.h +++ b/source/include/rocprofiler-sdk/hsa/api_trace_version.h @@ -22,7 +22,7 @@ #pragma once -#include +#include #if defined(__cplusplus) # include // safe to include from C++ diff --git a/source/include/rocprofiler-sdk/hsa/core_api_id.h b/source/include/rocprofiler-sdk/hsa/core_api_id.h index 134b5d55da..6262e28c26 100644 --- a/source/include/rocprofiler-sdk/hsa/core_api_id.h +++ b/source/include/rocprofiler-sdk/hsa/core_api_id.h @@ -22,12 +22,12 @@ #pragma once -#include +#include /** * @brief ROCProfiler enumeration of HSA Core API tracing operations */ -typedef enum // NOLINT(performance-enum-size) +typedef enum rocprofiler_hsa_core_api_id_t // NOLINT(performance-enum-size) { ROCPROFILER_HSA_CORE_API_ID_NONE = -1, ROCPROFILER_HSA_CORE_API_ID_hsa_init = 0, diff --git a/source/include/rocprofiler-sdk/hsa/finalize_ext_api_id.h b/source/include/rocprofiler-sdk/hsa/finalize_ext_api_id.h index 3c5acce424..b4fd0b92b3 100644 --- a/source/include/rocprofiler-sdk/hsa/finalize_ext_api_id.h +++ b/source/include/rocprofiler-sdk/hsa/finalize_ext_api_id.h @@ -22,12 +22,12 @@ #pragma once -#include +#include /** * @brief ROCProfiler enumeration of HSA Image Extended API tracing operations */ -typedef enum // NOLINT(performance-enum-size) +typedef enum rocprofiler_hsa_finalize_ext_api_id_t // NOLINT(performance-enum-size) { ROCPROFILER_HSA_FINALIZE_EXT_API_ID_NONE = -1, ROCPROFILER_HSA_FINALIZE_EXT_API_ID_hsa_ext_program_create = 0, diff --git a/source/include/rocprofiler-sdk/hsa/image_ext_api_id.h b/source/include/rocprofiler-sdk/hsa/image_ext_api_id.h index 5eb61a6600..d9cf9ccc3e 100644 --- a/source/include/rocprofiler-sdk/hsa/image_ext_api_id.h +++ b/source/include/rocprofiler-sdk/hsa/image_ext_api_id.h @@ -22,12 +22,12 @@ #pragma once -#include +#include /** * @brief ROCProfiler enumeration of HSA Image Extended API tracing operations */ -typedef enum // NOLINT(performance-enum-size) +typedef enum rocprofiler_hsa_image_ext_api_id_t // NOLINT(performance-enum-size) { ROCPROFILER_HSA_IMAGE_EXT_API_ID_NONE = -1, ROCPROFILER_HSA_IMAGE_EXT_API_ID_hsa_ext_image_get_capability, diff --git a/source/include/rocprofiler-sdk/hsa/scratch_memory_args.h b/source/include/rocprofiler-sdk/hsa/scratch_memory_args.h index f1e49f3a61..137131357d 100644 --- a/source/include/rocprofiler-sdk/hsa/scratch_memory_args.h +++ b/source/include/rocprofiler-sdk/hsa/scratch_memory_args.h @@ -23,11 +23,9 @@ #pragma once #include -#include -#include -#include -#include +#include +#include ROCPROFILER_EXTERN_C_INIT diff --git a/source/include/rocprofiler-sdk/hsa/scratch_memory_id.h b/source/include/rocprofiler-sdk/hsa/scratch_memory_id.h index 9a7de11dc4..c9df0db9cf 100644 --- a/source/include/rocprofiler-sdk/hsa/scratch_memory_id.h +++ b/source/include/rocprofiler-sdk/hsa/scratch_memory_id.h @@ -28,7 +28,7 @@ * @brief Allocation flags for @see rocprofiler_buffer_tracing_scratch_memory_record_t */ // NOLINTNEXTLINE(performance-enum-size) -typedef enum +typedef enum rocprofiler_scratch_alloc_flag_t { ROCPROFILER_SCRATCH_ALLOC_FLAG_NONE = 0, ROCPROFILER_SCRATCH_ALLOC_FLAG_USE_ONCE = diff --git a/source/include/rocprofiler-sdk/hsa/table_id.h b/source/include/rocprofiler-sdk/hsa/table_id.h index 59fdec1013..4205ce63b2 100644 --- a/source/include/rocprofiler-sdk/hsa/table_id.h +++ b/source/include/rocprofiler-sdk/hsa/table_id.h @@ -23,7 +23,7 @@ #pragma once // NOLINTNEXTLINE(performance-enum-size) -typedef enum +typedef enum rocprofiler_hsa_table_id_t { ROCPROFILER_HSA_TABLE_ID_NONE = -1, ROCPROFILER_HSA_TABLE_ID_Core = 0, diff --git a/source/include/rocprofiler-sdk/intercept_table.h b/source/include/rocprofiler-sdk/intercept_table.h index 5f3462b888..3ff70f88bd 100644 --- a/source/include/rocprofiler-sdk/intercept_table.h +++ b/source/include/rocprofiler-sdk/intercept_table.h @@ -49,17 +49,18 @@ ROCPROFILER_EXTERN_C_INIT */ /** - * @brief Callback type when a new runtime library is loaded. @see + * @brief (experimental) Callback type when a new runtime library is loaded. @see * rocprofiler_at_intercept_table_registration * @param [in] type Type of API table * @param [in] lib_version Major, minor, and patch version of library encoded into single number - * similar to @ref ROCPROFILER_VERSION + * similar to ::ROCPROFILER_VERSION * @param [in] lib_instance The number of times this runtime library has been registered previously * @param [in] tables An array of pointers to the API tables * @param [in] num_tables The size of the array of pointers to the API tables - * @param [in] user_data The pointer to the data provided to @ref - * rocprofiler_at_intercept_table_registration + * @param [in] user_data The pointer to the data provided to + * ::rocprofiler_at_intercept_table_registration */ +ROCPROFILER_SDK_EXPERIMENTAL typedef void (*rocprofiler_intercept_library_cb_t)(rocprofiler_intercept_table_t type, uint64_t lib_version, uint64_t lib_instance, @@ -68,10 +69,10 @@ typedef void (*rocprofiler_intercept_library_cb_t)(rocprofiler_intercept_table_t void* user_data); /** - * @brief Invoke this function to receive callbacks when a ROCm library registers its API - * intercept table with rocprofiler. Use the @ref rocprofiler_intercept_table_t enumeration for - * specifying which raw API tables the tool would like to have access to. E.g. including @ref - * ROCPROFILER_HSA_TABLE in the @ref rocprofiler_at_intercept_table_registration function call + * @brief (experimental) Invoke this function to receive callbacks when a ROCm library registers its + * API intercept table with rocprofiler. Use the ::rocprofiler_intercept_table_t enumeration for + * specifying which raw API tables the tool would like to have access to. E.g. including + * ::ROCPROFILER_HSA_TABLE in the ::rocprofiler_at_intercept_table_registration function call * communicates to rocprofiler that, when rocprofiler receives a `HsaApiTable` instance, the tool * would like rocprofiler to provide it access too. * @@ -79,23 +80,23 @@ typedef void (*rocprofiler_intercept_library_cb_t)(rocprofiler_intercept_table_t * invocation of one of their public API functions), these runtimes will provide a table of function * pointers to the rocprofiler library via the rocprofiler-register library if the * `rocprofiler_configure` symbol is visible in the application's symbol table. The vast majority of - * tools will want to use the @ref CALLBACK_TRACING_SERVICE to trace these runtime APIs, however, + * tools will want to use the ::CALLBACK_TRACING_SERVICE to trace these runtime APIs, however, * some tools may want or require installing their own intercept functions in lieu of receiving - * these callbacks and those tools should use the @ref rocprofiler_at_intercept_table_registration + * these callbacks and those tools should use the ::rocprofiler_at_intercept_table_registration * to install their intercept functions. There are no restrictions to where or how early this * function can be invoked but it will return ::ROCPROFILER_STATUS_ERROR_CONFIGURATION_LOCKED if it * is invoked after rocprofiler has requested all the tool configurations. Thus, it is highly - * recommended to invoke this function within the @ref rocprofiler_configure function or the - * callback passed to the @ref rocprofiler_force_configure function -- the reason for this - * recommendation is that if @ref rocprofiler_at_intercept_table_registration is invoked in one of + * recommended to invoke this function within the ::rocprofiler_configure function or the + * callback passed to the ::rocprofiler_force_configure function -- the reason for this + * recommendation is that if ::rocprofiler_at_intercept_table_registration is invoked in one of * these locations, rocprofiler can guarantee that the tool will be passed the API table because, at * the first instance of a runtime registering it's API table, rocprofiler will ensure that, in the - * case of the former, rocprofiler will invoke all of the @ref rocprofiler_configure symbols that + * case of the former, rocprofiler will invoke all of the ::rocprofiler_configure symbols that * are visible before checking the list of tools which want to receive the API tables and, in the - * case of the latter, @ref rocprofiler_force_configure will fail with error code @ref - * ROCPROFILER_STATUS_ERROR_CONFIGURATION_LOCKED if a runtime has already been registered (and, - * therefore, already scanned and invoked the visible @ref rocprofiler_configure symbols and - * completed the tool initialization). If @ref rocprofiler_at_intercept_table_registration is + * case of the latter, ::rocprofiler_force_configure will fail with error code + * ::ROCPROFILER_STATUS_ERROR_CONFIGURATION_LOCKED if a runtime has already been registered (and, + * therefore, already scanned and invoked the visible ::rocprofiler_configure symbols and + * completed the tool initialization). If ::rocprofiler_at_intercept_table_registration is * invoked outside of these recommended places, even if it is done before the `main` function starts * (e.g. in a library init/constructor function), it is possible that another library, such as * ROCm-aware MPI, caused the HIP and HSA runtime libraries to be initialized when that library was @@ -212,8 +213,9 @@ typedef void (*rocprofiler_intercept_library_cb_t)(rocprofiler_intercept_table_t * @endcode * * @example intercept_table/client.cpp - * Example demonstrating @ref rocprofiler_at_intercept_table_registration usage + * Example demonstrating ::rocprofiler_at_intercept_table_registration usage */ +ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_status_t rocprofiler_at_intercept_table_registration(rocprofiler_intercept_library_cb_t callback, int libs, diff --git a/source/include/rocprofiler-sdk/internal_threading.h b/source/include/rocprofiler-sdk/internal_threading.h index 14314efbcf..97eef290f2 100644 --- a/source/include/rocprofiler-sdk/internal_threading.h +++ b/source/include/rocprofiler-sdk/internal_threading.h @@ -33,23 +33,25 @@ ROCPROFILER_EXTERN_C_INIT * * @{ * @example api_buffered_tracing/client.cpp - * Example demonstrating @ref BUFFER_TRACING_SERVICE that includes usage of @ref - * rocprofiler_at_internal_thread_create, @ref rocprofiler_create_callback_thread, and @ref - * rocprofiler_assign_callback_thread. + * Example demonstrating ::BUFFER_TRACING_SERVICE that includes usage of + * ::rocprofiler_at_internal_thread_create, ::rocprofiler_create_callback_thread, and + * ::rocprofiler_assign_callback_thread. */ /** - * @brief Callback type before and after internal thread creation. @see + * @brief (experimental) Callback type before and after internal thread creation. @see * rocprofiler_at_internal_thread_create * */ +ROCPROFILER_SDK_EXPERIMENTAL typedef void (*rocprofiler_internal_thread_library_cb_t)(rocprofiler_runtime_library_t, void*); /** - * @brief Invoke this function to receive callbacks before and after the creation of an internal - * thread by a library which as invoked on the thread which is creating the internal thread(s). + * @brief (experimental) Invoke this function to receive callbacks before and after the creation of + * an internal thread by a library which as invoked on the thread which is creating the internal + * thread(s). * - * Use the @ref rocprofiler_runtime_library_t enumeration for specifying which libraries you want + * Use the ::rocprofiler_runtime_library_t enumeration for specifying which libraries you want * callbacks before and after the library creates an internal thread. These callbacks will be * invoked on the thread that is about to create the new thread (not on the newly created thread). * In thread-aware tools that wrap pthread_create, this can be used to disable the wrapper before @@ -74,6 +76,7 @@ typedef void (*rocprofiler_internal_thread_library_cb_t)(rocprofiler_runtime_lib * @retval ::ROCPROFILER_STATUS_SUCCESS There are currently no conditions which result in any other * value, even if internal threads have already been created */ +ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_status_t rocprofiler_at_internal_thread_create(rocprofiler_internal_thread_library_cb_t precreate, rocprofiler_internal_thread_library_cb_t postcreate, @@ -81,38 +84,40 @@ rocprofiler_at_internal_thread_create(rocprofiler_internal_thread_library_cb_t p void* data) ROCPROFILER_API; /** - * @brief opaque handle to an internal thread identifier which delivers callbacks for buffers + * @brief (experimental) opaque handle to an internal thread identifier which delivers callbacks for + * buffers * @see rocprofiler_create_callback_thread */ -typedef struct +typedef struct ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_callback_thread_t { uint64_t handle; } rocprofiler_callback_thread_t; /** - * @brief Create a handle to a unique thread (created by rocprofiler) which, when associated with a - * particular buffer, will guarantee those buffered results always get delivered on the same thread. - * This is useful to prevent/control thread-safety issues and/or enable multithreaded processing of - * buffers with non-overlapping data + * @brief (experimental) Create a handle to a unique thread (created by rocprofiler) which, when + * associated with a particular buffer, will guarantee those buffered results always get delivered + * on the same thread. This is useful to prevent/control thread-safety issues and/or enable + * multithreaded processing of buffers with non-overlapping data * - * @param [in] cb_thread_id User-provided pointer to a @ref rocprofiler_callback_thread_t + * @param [in] cb_thread_id User-provided pointer to a ::rocprofiler_callback_thread_t * @return ::rocprofiler_status_t * @retval ::ROCPROFILER_STATUS_SUCCESS Successful thread creation * @retval ::ROCPROFILER_STATUS_ERROR_CONFIGURATION_LOCKED Thread creation is no longer available * post-initialization * @retval ::ROCPROFILER_STATUS_ERROR Failed to create thread */ +ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_status_t rocprofiler_create_callback_thread(rocprofiler_callback_thread_t* cb_thread_id) ROCPROFILER_API ROCPROFILER_NONNULL(1); /** - * @brief By default, all buffered results are delivered on the same thread. Using @ref - * rocprofiler_create_callback_thread, one or more buffers can be assigned to deliever their results - * on a unique, dedicated thread. + * @brief (experimental) By default, all buffered results are delivered on the same thread. Using + * ::rocprofiler_create_callback_thread, one or more buffers can be assigned to deliever their + * results on a unique, dedicated thread. * * @param [in] buffer_id Buffer identifier - * @param [in] cb_thread_id Callback thread identifier via @ref rocprofiler_create_callback_thread + * @param [in] cb_thread_id Callback thread identifier via ::rocprofiler_create_callback_thread * @return ::rocprofiler_status_t * @retval ::ROCPROFILER_STATUS_SUCCESS Successful assignment of the delivery thread for the given * buffer @@ -123,6 +128,7 @@ rocprofiler_create_callback_thread(rocprofiler_callback_thread_t* cb_thread_id) * @retval ::ROCPROFILER_STATUS_ERROR_BUFFER_NOT_FOUND Buffer identifier did not match any of the * buffers registered with rocprofiler */ +ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_status_t rocprofiler_assign_callback_thread(rocprofiler_buffer_id_t buffer_id, rocprofiler_callback_thread_t cb_thread_id) ROCPROFILER_API; diff --git a/source/include/rocprofiler-sdk/kfd/page_migration_args.h b/source/include/rocprofiler-sdk/kfd/page_migration_args.h index ad13942ad7..dfc9e8f424 100644 --- a/source/include/rocprofiler-sdk/kfd/page_migration_args.h +++ b/source/include/rocprofiler-sdk/kfd/page_migration_args.h @@ -94,7 +94,7 @@ typedef struct rocprofiler_page_migration_dropped_event_t uint32_t dropped_events_count; } rocprofiler_page_migration_dropped_event_t; -typedef union +typedef union rocprofiler_page_migration_args_t { rocprofiler_page_migration_none_t none; rocprofiler_page_migration_page_migrate_start_t page_migrate_start; diff --git a/source/include/rocprofiler-sdk/kfd/page_migration_id.h b/source/include/rocprofiler-sdk/kfd/page_migration_id.h index 6ba74f6cb4..53f3b480b4 100644 --- a/source/include/rocprofiler-sdk/kfd/page_migration_id.h +++ b/source/include/rocprofiler-sdk/kfd/page_migration_id.h @@ -26,7 +26,6 @@ #include #include #include -#include #include @@ -36,7 +35,7 @@ ROCPROFILER_EXTERN_C_INIT * @brief Page migration triggers * */ -typedef enum +typedef enum rocprofiler_page_migration_trigger_t { ROCPROFILER_PAGE_MIGRATION_TRIGGER_NONE = -1, ROCPROFILER_PAGE_MIGRATION_TRIGGER_PREFETCH, ///< Migration triggered by a prefetch @@ -51,7 +50,7 @@ typedef enum * @brief Page migration triggers causing the queue to suspend * */ -typedef enum +typedef enum rocprofiler_page_migration_queue_suspend_trigger_t { ROCPROFILER_PAGE_MIGRATION_QUEUE_SUSPEND_TRIGGER_NONE = -1, ROCPROFILER_PAGE_MIGRATION_QUEUE_SUSPEND_TRIGGER_SVM, @@ -71,7 +70,7 @@ typedef enum * @brief Page migration triggers causing an unmap from the GPU * */ -typedef enum +typedef enum rocprofiler_page_migration_unmap_from_gpu_trigger_t { ROCPROFILER_PAGE_MIGRATION_UNMAP_FROM_GPU_TRIGGER_NONE = -1, ROCPROFILER_PAGE_MIGRATION_UNMAP_FROM_GPU_TRIGGER_MMU_NOTIFY, diff --git a/source/include/rocprofiler-sdk/marker/api_args.h b/source/include/rocprofiler-sdk/marker/api_args.h index 009c68755b..aff3e10c0d 100644 --- a/source/include/rocprofiler-sdk/marker/api_args.h +++ b/source/include/rocprofiler-sdk/marker/api_args.h @@ -23,7 +23,6 @@ #pragma once #include -#include #include #include diff --git a/source/include/rocprofiler-sdk/marker/api_id.h b/source/include/rocprofiler-sdk/marker/api_id.h index 273f5e2142..2b30f19658 100644 --- a/source/include/rocprofiler-sdk/marker/api_id.h +++ b/source/include/rocprofiler-sdk/marker/api_id.h @@ -25,7 +25,7 @@ /** * @brief ROCProfiler enumeration of Marker (ROCTx) API tracing operations */ -typedef enum // NOLINT(performance-enum-size) +typedef enum rocprofiler_marker_core_api_id_t // NOLINT(performance-enum-size) { ROCPROFILER_MARKER_CORE_API_ID_NONE = -1, ROCPROFILER_MARKER_CORE_API_ID_roctxMarkA = 0, @@ -37,7 +37,7 @@ typedef enum // NOLINT(performance-enum-size) ROCPROFILER_MARKER_CORE_API_ID_LAST, } rocprofiler_marker_core_api_id_t; -typedef enum // NOLINT(performance-enum-size) +typedef enum rocprofiler_marker_control_api_id_t // NOLINT(performance-enum-size) { ROCPROFILER_MARKER_CONTROL_API_ID_NONE = -1, ROCPROFILER_MARKER_CONTROL_API_ID_roctxProfilerPause = 0, @@ -45,7 +45,7 @@ typedef enum // NOLINT(performance-enum-size) ROCPROFILER_MARKER_CONTROL_API_ID_LAST, } rocprofiler_marker_control_api_id_t; -typedef enum // NOLINT(performance-enum-size) +typedef enum rocprofiler_marker_name_api_id_t // NOLINT(performance-enum-size) { ROCPROFILER_MARKER_NAME_API_ID_NONE = -1, ROCPROFILER_MARKER_NAME_API_ID_roctxNameOsThread = 0, diff --git a/source/include/rocprofiler-sdk/marker/table_id.h b/source/include/rocprofiler-sdk/marker/table_id.h index 4eaf1ab118..07cf3e406e 100644 --- a/source/include/rocprofiler-sdk/marker/table_id.h +++ b/source/include/rocprofiler-sdk/marker/table_id.h @@ -23,7 +23,7 @@ #pragma once // NOLINTNEXTLINE(performance-enum-size) -typedef enum +typedef enum rocprofiler_marker_table_id_t { ROCPROFILER_MARKER_TABLE_ID_NONE = -1, ROCPROFILER_MARKER_TABLE_ID_RoctxCore = 0, diff --git a/source/include/rocprofiler-sdk/ompt.h b/source/include/rocprofiler-sdk/ompt.h index 5fae859e61..9211a46f96 100644 --- a/source/include/rocprofiler-sdk/ompt.h +++ b/source/include/rocprofiler-sdk/ompt.h @@ -39,12 +39,40 @@ ROCPROFILER_EXTERN_C_INIT +/** + * @brief (experimental) Query whether rocprofiler-sdk OMPT implementation has been initialized by + * OpenMP runtime. + * + * @param [out] status Set to 0 if rocprofiler OMPT has not been initialized. Otherwise, set to 1. + * @return ::rocprofiler_status_t + * @retval ::ROCPROFILER_STATUS_SUCCESS Always returns this value + */ +ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_status_t rocprofiler_ompt_is_initialized(int* status) ROCPROFILER_API ROCPROFILER_NONNULL(1); +/** + * @brief (experimental) Query whether rocprofiler-sdk OMPT implementation has invoked ompt_finalize + * function. + * + * @param [out] status Set to 0 if rocprofiler OMPT has not been finalized. Otherwise, set to 1. + * @return ::rocprofiler_status_t + * @retval ::ROCPROFILER_STATUS_SUCCESS Always returns this value + */ +ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_status_t rocprofiler_ompt_is_finalized(int* status) ROCPROFILER_API ROCPROFILER_NONNULL(1); +/** + * @brief (experimental) If a tool which contains a "ompt_start_tool" function which is invoked by + * the OpenMP runtime but the tool wishes to defer to rocprofiler-sdk to be the OMPT tool, it should + * invoke this function from it's `ompt_start_tool` implementation. + * + * @param [in] omp_version Refer to OpenMP OMPT docs for more information + * @param [in] runtime_version Refer to OpenMP OMPT docs for more information + * @return ompt_start_tool_result_t* + */ +ROCPROFILER_SDK_EXPERIMENTAL ompt_start_tool_result_t* rocprofiler_ompt_start_tool(unsigned int omp_version, const char* runtime_version) ROCPROFILER_API; diff --git a/source/include/rocprofiler-sdk/ompt/api_args.h b/source/include/rocprofiler-sdk/ompt/api_args.h index 6677e4ef7f..ad0a61368b 100644 --- a/source/include/rocprofiler-sdk/ompt/api_args.h +++ b/source/include/rocprofiler-sdk/ompt/api_args.h @@ -23,7 +23,6 @@ #pragma once #include -#include #include diff --git a/source/include/rocprofiler-sdk/ompt/api_id.h b/source/include/rocprofiler-sdk/ompt/api_id.h index f78b54cd42..ab6079389e 100644 --- a/source/include/rocprofiler-sdk/ompt/api_id.h +++ b/source/include/rocprofiler-sdk/ompt/api_id.h @@ -26,7 +26,7 @@ * @brief ROCProfiler enumeration of OMPT (OpenMP tools) tracing operations * NOTE: These are callbacks into the ROCProfiler SDK from the vendor-provided OMPT implementation */ -typedef enum // NOLINT(performance-enum-size) +typedef enum rocprofiler_ompt_operation_t // NOLINT(performance-enum-size) { ROCPROFILER_OMPT_ID_NONE = -1, ROCPROFILER_OMPT_ID_thread_begin = 0, diff --git a/source/include/rocprofiler-sdk/pc_sampling.h b/source/include/rocprofiler-sdk/pc_sampling.h index 40d8293666..f23611ed52 100644 --- a/source/include/rocprofiler-sdk/pc_sampling.h +++ b/source/include/rocprofiler-sdk/pc_sampling.h @@ -35,7 +35,8 @@ ROCPROFILER_EXTERN_C_INIT */ /** - * @brief Function used to configure the PC sampling service on the GPU agent with @p agent_id. + * @brief (experimental) Function used to configure the PC sampling service on the GPU agent with @p + * agent_id. * * Prerequisites are the following: * - The client must create a context and supply its @p context_id. By using this context, @@ -114,6 +115,7 @@ ROCPROFILER_EXTERN_C_INIT * setup in the context * @retval ::ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT function invoked with an invalid argument */ +ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_status_t rocprofiler_configure_pc_sampling_service(rocprofiler_context_id_t context_id, rocprofiler_agent_id_t agent_id, @@ -124,9 +126,10 @@ rocprofiler_configure_pc_sampling_service(rocprofiler_context_id_t conte int flags) ROCPROFILER_API; /** - * @brief Enumeration describing values of flags of ::rocprofiler_pc_sampling_configuration_t. + * @brief (experimental) Enumeration describing values of flags of + * ::rocprofiler_pc_sampling_configuration_t. */ -typedef enum rocprofiler_pc_sampling_configuration_flags_t +typedef enum ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_pc_sampling_configuration_flags_t { ROCPROFILER_PC_SAMPLING_CONFIGURATION_FLAGS_NONE = 0, ROCPROFILER_PC_SAMPLING_CONFIGURATION_FLAGS_INTERVAL_POW2, @@ -137,9 +140,9 @@ typedef enum rocprofiler_pc_sampling_configuration_flags_t } rocprofiler_pc_sampling_configuration_flags_t; /** - * @brief PC sampling configuration supported by a GPU agent. + * @brief (experimental) PC sampling configuration supported by a GPU agent. */ -typedef struct +typedef struct ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_pc_sampling_configuration_t { uint64_t size; ///< Size of this struct rocprofiler_pc_sampling_method_t method; @@ -165,26 +168,27 @@ typedef struct } rocprofiler_pc_sampling_configuration_t; /** - * @brief Rocprofiler SDK's callback function to deliver the list of available PC + * @brief (experimental) Rocprofiler SDK's callback function to deliver the list of available PC * sampling configurations upon the call to the - * @ref rocprofiler_query_pc_sampling_agent_configurations. + * ::rocprofiler_query_pc_sampling_agent_configurations. * * @param[out] configs - The array of PC sampling configurations supported by the agent - * at the moment of invoking @ref rocprofiler_query_pc_sampling_agent_configurations. + * at the moment of invoking ::rocprofiler_query_pc_sampling_agent_configurations. * @param[out] num_config - The number of configurations contained in the underlying array * @p configs. * In case the GPU agent does not support PC sampling, the value is 0. * @param[in] user_data - client's private data passed via - * @ref rocprofiler_query_pc_sampling_agent_configurations + * ::rocprofiler_query_pc_sampling_agent_configurations * @return ::rocprofiler_status_t */ +ROCPROFILER_SDK_EXPERIMENTAL typedef rocprofiler_status_t (*rocprofiler_available_pc_sampling_configurations_cb_t)( const rocprofiler_pc_sampling_configuration_t* configs, size_t num_config, void* user_data); /** - * @brief Query PC Sampling Configuration. + * @brief (experimental) Query PC Sampling Configuration. * * Lists PC sampling configurations a GPU agent with @p agent_id supports at the moment * of invoking the function. Delivers configurations via @p cb. @@ -205,6 +209,7 @@ typedef rocprofiler_status_t (*rocprofiler_available_pc_sampling_configurations_ * @retval ::ROCPROFILER_STATUS_ERROR a general error caused by the amdgpu driver * @retval ::ROCPROFILER_STATUS_SUCCESS @p cb successfully finished */ +ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_status_t rocprofiler_query_pc_sampling_agent_configurations( rocprofiler_agent_id_t agent_id, @@ -212,10 +217,10 @@ rocprofiler_query_pc_sampling_agent_configurations( void* user_data) ROCPROFILER_API ROCPROFILER_NONNULL(2, 3); /** - * @brief Information about the GPU part where wave was executing + * @brief (experimental) Information about the GPU part where wave was executing * at the moment of sampling. */ -typedef struct rocprofiler_pc_sampling_hw_id_v0_t +typedef struct ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_pc_sampling_hw_id_v0_t { uint64_t chiplet : 6; ///< chiplet index (3 bits allocated by the ROCr runtime) uint64_t wave_id : 7; ///< wave slot index @@ -235,9 +240,9 @@ typedef struct rocprofiler_pc_sampling_hw_id_v0_t } rocprofiler_pc_sampling_hw_id_v0_t; /** - * @brief Sampled program counter. + * @brief (experimental) Sampled program counter. */ -typedef struct +typedef struct ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_pc_t { uint64_t code_object_id; uint64_t code_object_offset; @@ -260,9 +265,9 @@ typedef struct } rocprofiler_pc_t; /** - * @brief ROCProfiler Host-Trap PC Sampling Record. + * @brief (experimental) ROCProfiler Host-Trap PC Sampling Record. */ -typedef struct rocprofiler_pc_sampling_record_host_trap_v0_t +typedef struct ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_pc_sampling_record_host_trap_v0_t { uint64_t size; ///< Size of this struct rocprofiler_pc_sampling_hw_id_v0_t hw_id; ///< @see ::rocprofiler_pc_sampling_hw_id_0_t @@ -280,11 +285,11 @@ typedef struct rocprofiler_pc_sampling_record_host_trap_v0_t } rocprofiler_pc_sampling_record_host_trap_v0_t; /** - * @brief The header of the @ref rocprofiler_pc_sampling_record_stochastic_v0_t, indicating - * what fields of the @ref rocprofiler_pc_sampling_record_stochastic_v0_t instance are meaningful - * for the sample. + * @brief (experimental) The header of the ::rocprofiler_pc_sampling_record_stochastic_v0_t, + * indicating what fields of the ::rocprofiler_pc_sampling_record_stochastic_v0_t instance are + * meaningful for the sample. */ -typedef struct rocprofiler_pc_sampling_record_stochastic_header_t +typedef struct ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_pc_sampling_record_stochastic_header_t { uint8_t has_memory_counter : 1; ///< pc sample provides memory counters information ///< via ::rocprofiler_pc_sampling_memory_counters_t @@ -292,9 +297,9 @@ typedef struct rocprofiler_pc_sampling_record_stochastic_header_t } rocprofiler_pc_sampling_record_stochastic_header_t; /** - * @brief Enumeration describing type of sampled issued instruction. + * @brief (experimental) Enumeration describing type of sampled issued instruction. */ -typedef enum rocprofiler_pc_sampling_instruction_type_t +typedef enum ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_pc_sampling_instruction_type_t { ROCPROFILER_PC_SAMPLING_INSTRUCTION_TYPE_NONE = 0, ROCPROFILER_PC_SAMPLING_INSTRUCTION_TYPE_VALU, ///< vector ALU instruction @@ -322,9 +327,9 @@ typedef enum rocprofiler_pc_sampling_instruction_type_t } rocprofiler_pc_sampling_instruction_type_t; /** - * @brief Enumeration describing reason for not issuing an instruction. + * @brief (experimental) Enumeration describing reason for not issuing an instruction. */ -typedef enum rocprofiler_pc_sampling_instruction_not_issued_reason_t +typedef enum ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_pc_sampling_instruction_not_issued_reason_t { ROCPROFILER_PC_SAMPLING_INSTRUCTION_NOT_ISSUED_REASON_NONE = 0, ROCPROFILER_PC_SAMPLING_INSTRUCTION_NOT_ISSUED_REASON_NO_INSTRUCTION_AVAILABLE, @@ -354,10 +359,10 @@ typedef enum rocprofiler_pc_sampling_instruction_not_issued_reason_t } rocprofiler_pc_sampling_instruction_not_issued_reason_t; /** - * @brief Data provided by stochastic sampling hardware. + * @brief (experimental) Data provided by stochastic sampling hardware. * */ -typedef struct rocprofiler_pc_sampling_snapshot_v0_t +typedef struct ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_pc_sampling_snapshot_v0_t { uint32_t reason_not_issued : 4; uint32_t reserved0 : 1; ///< reserved for future use @@ -390,7 +395,7 @@ typedef struct rocprofiler_pc_sampling_snapshot_v0_t /// @var reason_not_issued /// @brief The reason for not issuing an instruction. The field takes one of the value defined - /// in @ref ::rocprofiler_pc_sampling_instruction_not_issued_reason_t + /// in ::rocprofiler_pc_sampling_instruction_not_issued_reason_t /// @var arb_state_stall_valu /// @brief VALU instruction was stalled when a sample was generated /// @var dual_issue_valu @@ -398,9 +403,9 @@ typedef struct rocprofiler_pc_sampling_snapshot_v0_t } rocprofiler_pc_sampling_snapshot_v0_t; /** - * @brief Counters of issued but not yet completed instructions. + * @brief (experimental) Counters of issued but not yet completed instructions. */ -typedef struct rocprofiler_pc_sampling_memory_counters_t +typedef struct ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_pc_sampling_memory_counters_t { uint32_t load_cnt : 6; uint32_t store_cnt : 6; @@ -425,9 +430,9 @@ typedef struct rocprofiler_pc_sampling_memory_counters_t } rocprofiler_pc_sampling_memory_counters_t; /** - * @brief ROCProfiler Stochastic PC Sampling Record. + * @brief (experimental) ROCProfiler Stochastic PC Sampling Record. */ -typedef struct rocprofiler_pc_sampling_record_stochastic_v0_t +typedef struct ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_pc_sampling_record_stochastic_v0_t { uint64_t size; ///< Size of this struct rocprofiler_pc_sampling_record_stochastic_header_t flags; @@ -482,33 +487,33 @@ typedef struct rocprofiler_pc_sampling_record_stochastic_v0_t } rocprofiler_pc_sampling_record_stochastic_v0_t; /** - * @brief Record representing an invalid PC Sampling Record. + * @brief (experimental) Record representing an invalid PC Sampling Record. */ -typedef struct rocprofiler_pc_sampling_record_invalid_t +typedef struct ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_pc_sampling_record_invalid_t { uint64_t size; ///< Size of the struct } rocprofiler_pc_sampling_record_invalid_t; /** - * @fn C compatible string representation of the PC sampling instruction type - * @brief Return the string encoding of @ref rocprofiler_pc_sampling_instruction_type_t value + * @brief (experimental) Return the string encoding of ::rocprofiler_pc_sampling_instruction_type_t + * value * @param [in] instruction_type instruction type enum value - * @return Will return a nullptr if invalid/unsupported @ref - * rocprofiler_pc_sampling_instruction_type_t value is provided. + * @return Will return a nullptr if invalid/unsupported ::rocprofiler_pc_sampling_instruction_type_t + * value is provided. */ +ROCPROFILER_SDK_EXPERIMENTAL const char* rocprofiler_get_pc_sampling_instruction_type_name( rocprofiler_pc_sampling_instruction_type_t instruction_type) ROCPROFILER_API; /** - * @fn C compatible string representation of reason for not issuing an instruciton - * @brief Return the string encoding of @ref rocprofiler_pc_sampling_instruction_not_issued_reason_t - * value + * @brief (experimental) Return the string encoding of + * ::rocprofiler_pc_sampling_instruction_not_issued_reason_t value * @param [in] not_issued_reason no issue reason enum value - * @return Will return a nullptr if invalid/unsupported @ref - * rocprofiler_pc_sampling_instruction_not_issued_reason_t value is provided. + * @return Will return a nullptr if invalid/unsupported + * ::rocprofiler_pc_sampling_instruction_not_issued_reason_t value is provided. */ -const char* +ROCPROFILER_SDK_EXPERIMENTAL const char* rocprofiler_get_pc_sampling_instruction_not_issued_reason_name( rocprofiler_pc_sampling_instruction_not_issued_reason_t not_issued_reason) ROCPROFILER_API; diff --git a/source/include/rocprofiler-sdk/rccl/api_args.h b/source/include/rocprofiler-sdk/rccl/api_args.h index a78d4f9c61..3af8d07a42 100644 --- a/source/include/rocprofiler-sdk/rccl/api_args.h +++ b/source/include/rocprofiler-sdk/rccl/api_args.h @@ -23,7 +23,6 @@ #pragma once #include -#include #if !defined(ROCPROFILER_SDK_USE_SYSTEM_RCCL) # if defined __has_include diff --git a/source/include/rocprofiler-sdk/rccl/api_id.h b/source/include/rocprofiler-sdk/rccl/api_id.h index fc3f2a577a..1c668ab355 100644 --- a/source/include/rocprofiler-sdk/rccl/api_id.h +++ b/source/include/rocprofiler-sdk/rccl/api_id.h @@ -24,12 +24,10 @@ #pragma once -#include - /** - * @brief ROCProfiler enumeration of HSA Core API tracing operations + * @brief ROCProfiler enumeration of RCCL API tracing operations */ -typedef enum // NOLINT(performance-enum-size) +typedef enum rocprofiler_rccl_api_id_t // NOLINT(performance-enum-size) { ROCPROFILER_RCCL_API_ID_NONE = -1, diff --git a/source/include/rocprofiler-sdk/rccl/table_id.h b/source/include/rocprofiler-sdk/rccl/table_id.h index b4fb303daf..d6b2b1e683 100644 --- a/source/include/rocprofiler-sdk/rccl/table_id.h +++ b/source/include/rocprofiler-sdk/rccl/table_id.h @@ -23,7 +23,7 @@ #pragma once // NOLINTNEXTLINE(performance-enum-size) -typedef enum +typedef enum rocprofiler_rccl_table_id_t { ROCPROFILER_RCCL_TABLE_ID_NONE = -1, ROCPROFILER_RCCL_TABLE_ID = 0, diff --git a/source/include/rocprofiler-sdk/registration.h b/source/include/rocprofiler-sdk/registration.h index 68be6f8ba4..f016b8c64e 100644 --- a/source/include/rocprofiler-sdk/registration.h +++ b/source/include/rocprofiler-sdk/registration.h @@ -35,17 +35,18 @@ ROCPROFILER_EXTERN_C_INIT */ /** - * @brief A client refers to an individual or entity engaged in the configuration of ROCprofiler - * services. e.g: any third party tool like PAPI or any internal tool (Omnitrace). A pointer to this - * data structure is provided to the client tool initialization function. The name member can be set - * by the client to assist with debugging (e.g. rocprofiler cannot start your context because there - * is a conflicting context started by `` -- at least that is the plan). The handle member is - * a unique identifer assigned by rocprofiler for the client and the client can store it and pass it - * to the @ref rocprofiler_client_finalize_t function to force finalization (i.e. deactivate all of - * it's contexts) for the client. + * @brief (experimental) A client refers to an individual or entity engaged in the configuration of + * ROCprofiler services. e.g: any third party tool like PAPI or any internal tool (Omnitrace). A + * pointer to this data structure is provided to the client tool initialization function. The name + * member can be set by the client to assist with debugging (e.g. rocprofiler cannot start your + * context because there is a conflicting context started by `` -- at least that is the plan). + * The handle member is a unique identifer assigned by rocprofiler for the client and the client can + * store it and pass it to the ::rocprofiler_client_finalize_t function to force finalization + * (i.e. deactivate all of it's contexts) for the client. */ -typedef struct rocprofiler_client_id_t +typedef struct ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_client_id_t { + size_t size; ///< size of this struct const char* name; ///< clients should set this value for debugging const uint32_t handle; ///< internal handle } rocprofiler_client_id_t; @@ -127,11 +128,11 @@ rocprofiler_is_finalized(int* status) ROCPROFILER_API ROCPROFILER_NONNULL(1); /** * @brief This is the special function that tools define to enable rocprofiler support. The tool * should return a pointer to - * @ref rocprofiler_tool_configure_result_t which will contain a function pointer to (1) an + * ::rocprofiler_tool_configure_result_t which will contain a function pointer to (1) an * initialization function where all the contexts are created, (2) a finalization function (if * necessary) which will be invoked when rocprofiler shutdown and, (3) a pointer to any data that - * the tool wants communicated between the @ref rocprofiler_tool_configure_result_t::initialize and - * @ref rocprofiler_tool_configure_result_t::finalize functions. If the user + * the tool wants communicated between the ::rocprofiler_tool_configure_result_t::initialize and + * ::rocprofiler_tool_configure_result_t::finalize functions. If the user * * @param [in] version The version of rocprofiler: `(10000 * major) + (100 * minor) + patch` * @param [in] runtime_version String descriptor of the rocprofiler version and other relevant info. @@ -227,7 +228,7 @@ rocprofiler_configure(uint32_t version, // want the symbol to be visible when the user includes the header for the prototype /** - * @brief Function pointer typedef for @ref rocprofiler_configure function + * @brief Function pointer typedef for ::rocprofiler_configure function * @param [in] version The version of rocprofiler: `(10000 * major) + (100 * minor) + patch` * @param [in] runtime_version String descriptor of the rocprofiler version and other relevant info. * @param [in] priority How many client tools were initialized before this client tool @@ -243,7 +244,7 @@ typedef rocprofiler_tool_configure_result_t* (*rocprofiler_configure_func_t)( * @brief Function for explicitly registering a configuration with rocprofiler. This can be invoked * before any ROCm runtimes (lazily) initialize and context(s) can be started before the runtimes * initialize. - * @param [in] configure_func Address of @ref rocprofiler_configure function. A null pointer is + * @param [in] configure_func Address of ::rocprofiler_configure function. A null pointer is * acceptable if the address is not known * @return ::rocprofiler_status_t * @retval ::ROCPROFILER_STATUS_SUCCESS Registration was successfully triggered. diff --git a/source/include/rocprofiler-sdk/rocdecode/CMakeLists.txt b/source/include/rocprofiler-sdk/rocdecode/CMakeLists.txt index f2528fec3a..8a1ca1b736 100644 --- a/source/include/rocprofiler-sdk/rocdecode/CMakeLists.txt +++ b/source/include/rocprofiler-sdk/rocdecode/CMakeLists.txt @@ -3,7 +3,7 @@ # Installation of public rocDecode headers # # -set(ROCPROFILER_ROCDECODE_HEADER_FILES api_args.h api_id.h table_id.h) +set(ROCPROFILER_ROCDECODE_HEADER_FILES api_args.h api_id.h api_trace.h table_id.h) install( FILES ${ROCPROFILER_ROCDECODE_HEADER_FILES} diff --git a/source/include/rocprofiler-sdk/rocdecode/api_args.h b/source/include/rocprofiler-sdk/rocdecode/api_args.h index 1dfae53828..13e0378d71 100644 --- a/source/include/rocprofiler-sdk/rocdecode/api_args.h +++ b/source/include/rocprofiler-sdk/rocdecode/api_args.h @@ -23,7 +23,7 @@ #pragma once #include -#include +#include #include diff --git a/source/include/rocprofiler-sdk/rocdecode/api_id.h b/source/include/rocprofiler-sdk/rocdecode/api_id.h index af8e05b32f..a71b631bd5 100644 --- a/source/include/rocprofiler-sdk/rocdecode/api_id.h +++ b/source/include/rocprofiler-sdk/rocdecode/api_id.h @@ -24,12 +24,12 @@ #pragma once -#include +#include /** * @brief ROCProfiler enumeration of rocDecode API tracing operations */ -typedef enum // NOLINT(performance-enum-size) +typedef enum rocprofiler_rocdecode_api_id_t // NOLINT(performance-enum-size) { ROCPROFILER_ROCDECODE_API_ID_NONE = -1, diff --git a/source/include/rocprofiler-sdk/rocdecode/api_trace.h b/source/include/rocprofiler-sdk/rocdecode/api_trace.h new file mode 100644 index 0000000000..f8feb012db --- /dev/null +++ b/source/include/rocprofiler-sdk/rocdecode/api_trace.h @@ -0,0 +1,31 @@ +// MIT License +// +// Copyright (c) 2023-2025 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 + +#if ROCPROFILER_SDK_USE_SYSTEM_ROCDECODE > 0 +# include +#else +# include +#endif diff --git a/source/include/rocprofiler-sdk/rocdecode/table_id.h b/source/include/rocprofiler-sdk/rocdecode/table_id.h index e2c6d1a129..4aa2d1004c 100644 --- a/source/include/rocprofiler-sdk/rocdecode/table_id.h +++ b/source/include/rocprofiler-sdk/rocdecode/table_id.h @@ -23,7 +23,7 @@ #pragma once // NOLINTNEXTLINE(performance-enum-size) -typedef enum +typedef enum rocprofiler_rocdecode_table_id_t { ROCPROFILER_ROCDECODE_TABLE_ID_NONE = -1, ROCPROFILER_ROCDECODE_TABLE_ID_CORE = 0, diff --git a/source/include/rocprofiler-sdk/rocjpeg/api_args.h b/source/include/rocprofiler-sdk/rocjpeg/api_args.h index 20bd12584b..be677e0a15 100644 --- a/source/include/rocprofiler-sdk/rocjpeg/api_args.h +++ b/source/include/rocprofiler-sdk/rocjpeg/api_args.h @@ -23,7 +23,6 @@ #pragma once #include -#include #include diff --git a/source/include/rocprofiler-sdk/rocjpeg/api_id.h b/source/include/rocprofiler-sdk/rocjpeg/api_id.h index f653789912..20b9d5fc45 100644 --- a/source/include/rocprofiler-sdk/rocjpeg/api_id.h +++ b/source/include/rocprofiler-sdk/rocjpeg/api_id.h @@ -22,12 +22,10 @@ #pragma once -#include - /** * @brief ROCProfiler enumeration of rocJPEG API tracing operations */ -typedef enum // NOLINT(performance-enum-size) +typedef enum rocprofiler_rocjpeg_api_id_t // NOLINT(performance-enum-size) { ROCPROFILER_ROCJPEG_API_ID_NONE = -1, diff --git a/source/include/rocprofiler-sdk/rocjpeg/table_id.h b/source/include/rocprofiler-sdk/rocjpeg/table_id.h index 6f07ccd164..d6f2fe2097 100644 --- a/source/include/rocprofiler-sdk/rocjpeg/table_id.h +++ b/source/include/rocprofiler-sdk/rocjpeg/table_id.h @@ -22,8 +22,7 @@ #pragma once -// NOLINTNEXTLINE(performance-enum-size) -typedef enum +typedef enum rocprofiler_rocjpeg_table_id_t // NOLINT(performance-enum-size) { ROCPROFILER_ROCJPEG_TABLE_ID_NONE = -1, ROCPROFILER_ROCJPEG_TABLE_ID_CORE = 0, diff --git a/source/include/rocprofiler-sdk/rocprofiler.h b/source/include/rocprofiler-sdk/rocprofiler.h index baa3e4b169..50ab8f7d0c 100644 --- a/source/include/rocprofiler-sdk/rocprofiler.h +++ b/source/include/rocprofiler-sdk/rocprofiler.h @@ -77,9 +77,9 @@ ROCPROFILER_EXTERN_C_INIT * patch) * @brief Query the version of the installed library. * - * Return the version of the installed library. This can be used to check if - * it is compatible with this interface version. This function can be used - * even when the library is not initialized. + * Returns the version of the rocprofiler-sdk library loaded at runtime. This can be used to check + * if the runtime version is equal to or compatible with the version of rocprofiler-sdk used during + * compilation time. This function can be invoked before tool initialization. * * @param [out] major The major version number is stored if non-NULL. * @param [out] minor The minor version number is stored if non-NULL. @@ -88,6 +88,19 @@ ROCPROFILER_EXTERN_C_INIT rocprofiler_status_t rocprofiler_get_version(uint32_t* major, uint32_t* minor, uint32_t* patch) ROCPROFILER_API; +/** + * @brief Simplified alternative to ::rocprofiler_get_version + * + * Returns the version of the rocprofiler-sdk library loaded at runtime. This can be used to check + * if the runtime version is equal to or compatible with the version of rocprofiler-sdk used during + * compilation time. This function can be invoked before tool initialization. + * + * @param [out] info Pointer to version triplet struct which will be populated by the function call. + */ +rocprofiler_status_t +rocprofiler_get_version_triplet(rocprofiler_version_triplet_t* info) ROCPROFILER_API + ROCPROFILER_NONNULL(1); + ROCPROFILER_EXTERN_C_FINI /** @} */ @@ -97,6 +110,7 @@ ROCPROFILER_EXTERN_C_FINI #include "rocprofiler-sdk/buffer_tracing.h" #include "rocprofiler-sdk/callback_tracing.h" #include "rocprofiler-sdk/context.h" +#include "rocprofiler-sdk/counter_config.h" #include "rocprofiler-sdk/counters.h" #include "rocprofiler-sdk/device_counting_service.h" #include "rocprofiler-sdk/dispatch_counting_service.h" @@ -106,10 +120,16 @@ ROCPROFILER_EXTERN_C_FINI #include "rocprofiler-sdk/intercept_table.h" #include "rocprofiler-sdk/internal_threading.h" #include "rocprofiler-sdk/marker.h" +#include "rocprofiler-sdk/ompt.h" #include "rocprofiler-sdk/pc_sampling.h" -#include "rocprofiler-sdk/profile_config.h" +#include "rocprofiler-sdk/rccl.h" +#include "rocprofiler-sdk/rocdecode.h" +#include "rocprofiler-sdk/rocjpeg.h" // #include "rocprofiler-sdk/spm.h" +// subject to removal +#include "rocprofiler-sdk/deprecated/profile_config.h" + ROCPROFILER_EXTERN_C_INIT /** diff --git a/source/include/rocprofiler-sdk/spm.h b/source/include/rocprofiler-sdk/spm.h index 0699dabb0e..85887bed72 100644 --- a/source/include/rocprofiler-sdk/spm.h +++ b/source/include/rocprofiler-sdk/spm.h @@ -39,14 +39,15 @@ ROCPROFILER_EXTERN_C_INIT * * @param [in] context_id * @param [in] buffer_id - * @param [in] profile_config + * @param [in] counter_config * @param [in] interval * @return ::rocprofiler_status_t */ +ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_status_t rocprofiler_configure_spm_service(rocprofiler_context_id_t context_id, rocprofiler_buffer_id_t buffer_id, - rocprofiler_profile_config_id_t profile_config, + rocprofiler_counter_config_id_t counter_config, uint64_t interval) ROCPROFILER_API; /** @} */ diff --git a/source/include/rocprofiler-sdk/version.h.in b/source/include/rocprofiler-sdk/version.h.in index 85d4f0531b..db6011ad34 100644 --- a/source/include/rocprofiler-sdk/version.h.in +++ b/source/include/rocprofiler-sdk/version.h.in @@ -108,32 +108,11 @@ // compiler information #define ROCPROFILER_COMPILER_ID "@CMAKE_CXX_COMPILER_ID@" #define ROCPROFILER_COMPILER_VERSION "@CMAKE_CXX_COMPILER_VERSION@" - -#define ROCPROFILER_HSA_RUNTIME_VERSION_MAJOR @HSA_RUNTIME_VERSION_MAJOR@ -#define ROCPROFILER_HSA_RUNTIME_VERSION_MINOR @HSA_RUNTIME_VERSION_MINOR@ - -#cmakedefine ROCPROFILER_HSA_API_TABLE_MAJOR_VERSION @ROCPROFILER_HSA_API_TABLE_MAJOR_VERSION@ -#cmakedefine ROCPROFILER_HSA_CORE_API_TABLE_MAJOR_VERSION @ROCPROFILER_HSA_CORE_API_TABLE_MAJOR_VERSION@ -#cmakedefine ROCPROFILER_HSA_AMD_EXT_API_TABLE_MAJOR_VERSION @ROCPROFILER_HSA_AMD_EXT_API_TABLE_MAJOR_VERSION@ -#cmakedefine ROCPROFILER_HSA_FINALIZER_API_TABLE_MAJOR_VERSION @ROCPROFILER_HSA_FINALIZER_API_TABLE_MAJOR_VERSION@ -#cmakedefine ROCPROFILER_HSA_IMAGE_API_TABLE_MAJOR_VERSION @ROCPROFILER_HSA_IMAGE_API_TABLE_MAJOR_VERSION@ -#cmakedefine ROCPROFILER_HSA_AQLPROFILE_API_TABLE_MAJOR_VERSION @ROCPROFILER_HSA_AQLPROFILE_API_TABLE_MAJOR_VERSION@ -#cmakedefine ROCPROFILER_HSA_TOOLS_API_TABLE_MAJOR_VERSION @ROCPROFILER_HSA_TOOLS_API_TABLE_MAJOR_VERSION@ - -#cmakedefine ROCPROFILER_HSA_API_TABLE_STEP_VERSION @ROCPROFILER_HSA_API_TABLE_STEP_VERSION@ -#cmakedefine ROCPROFILER_HSA_CORE_API_TABLE_STEP_VERSION @ROCPROFILER_HSA_CORE_API_TABLE_STEP_VERSION@ -#cmakedefine ROCPROFILER_HSA_AMD_EXT_API_TABLE_STEP_VERSION @ROCPROFILER_HSA_AMD_EXT_API_TABLE_STEP_VERSION@ -#cmakedefine ROCPROFILER_HSA_FINALIZER_API_TABLE_STEP_VERSION @ROCPROFILER_HSA_FINALIZER_API_TABLE_STEP_VERSION@ -#cmakedefine ROCPROFILER_HSA_IMAGE_API_TABLE_STEP_VERSION @ROCPROFILER_HSA_IMAGE_API_TABLE_STEP_VERSION@ -#cmakedefine ROCPROFILER_HSA_AQLPROFILE_API_TABLE_STEP_VERSION @ROCPROFILER_HSA_AQLPROFILE_API_TABLE_STEP_VERSION@ -#cmakedefine ROCPROFILER_HSA_TOOLS_API_TABLE_STEP_VERSION @ROCPROFILER_HSA_TOOLS_API_TABLE_STEP_VERSION@ // clang-format on #define ROCPROFILER_VERSION \ ((10000 * ROCPROFILER_VERSION_MAJOR) + (100 * ROCPROFILER_VERSION_MINOR) + \ ROCPROFILER_VERSION_PATCH) -// latest hsa-runtime version supported -#define ROCPROFILER_HSA_RUNTIME_VERSION \ - ((10000 * ROCPROFILER_HSA_RUNTIME_VERSION_MAJOR) + \ - (100 * ROCPROFILER_HSA_RUNTIME_VERSION_MINOR)) +// include the external version info +#include "ext_version.h" diff --git a/source/lib/common/abi.hpp b/source/lib/common/abi.hpp index a3d1eda962..11fe153f8e 100644 --- a/source/lib/common/abi.hpp +++ b/source/lib/common/abi.hpp @@ -22,7 +22,7 @@ #pragma once -#include +#include #include "lib/common/defines.hpp" diff --git a/source/lib/output/counter_info.hpp b/source/lib/output/counter_info.hpp index 39cd5165b5..b79871c9f1 100644 --- a/source/lib/output/counter_info.hpp +++ b/source/lib/output/counter_info.hpp @@ -43,9 +43,9 @@ constexpr uint32_t lds_block_size = 128 * 4; using counter_dimension_id_vec_t = std::vector; using counter_dimension_info_vec_t = std::vector; -struct tool_counter_info : rocprofiler_counter_info_v0_t +struct tool_counter_info : rocprofiler_counter_info_v1_t { - using parent_type = rocprofiler_counter_info_v0_t; + using parent_type = rocprofiler_counter_info_v1_t; tool_counter_info(rocprofiler_agent_id_t _agent_id, parent_type _info, @@ -125,7 +125,7 @@ void save(ArchiveT& ar, const ::rocprofiler::tool::tool_counter_info& data) { SAVE_DATA_FIELD(agent_id); - cereal::save(ar, static_cast(data)); + cereal::save(ar, static_cast(data)); SAVE_DATA_FIELD(dimension_ids); } diff --git a/source/lib/output/metadata.cpp b/source/lib/output/metadata.cpp index 850b685f04..d1c333af5d 100644 --- a/source/lib/output/metadata.cpp +++ b/source/lib/output/metadata.cpp @@ -49,19 +49,6 @@ namespace tool namespace fs = common::filesystem; namespace { -rocprofiler_status_t -dimensions_info_callback(rocprofiler_counter_id_t /*id*/, - const rocprofiler_record_dimension_info_t* dim_info, - long unsigned int num_dims, - void* user_data) -{ - auto* dimensions_info = static_cast(user_data); - dimensions_info->reserve(num_dims); - for(size_t j = 0; j < num_dims; j++) - dimensions_info->emplace_back(dim_info[j]); - return ROCPROFILER_STATUS_SUCCESS; -} - rocprofiler_status_t query_pc_sampling_configuration(const rocprofiler_pc_sampling_configuration_t* configs, long unsigned int num_config, @@ -177,31 +164,17 @@ void metadata::init(inprocess) for(size_t i = 0; i < num_counters; ++i) { - auto _info = rocprofiler_counter_info_v0_t{}; + auto _info = rocprofiler_counter_info_v1_t{}; auto _dim_ids = std::vector{}; auto _dim_info = std::vector{}; ROCPROFILER_CHECK(rocprofiler_query_counter_info( counters[i], - ROCPROFILER_COUNTER_INFO_VERSION_0, - &static_cast(_info))); + ROCPROFILER_COUNTER_INFO_VERSION_1, + &static_cast(_info))); - if(auto _itr_dim_stat = rocprofiler_iterate_counter_dimensions( - counters[i], dimensions_info_callback, &_dim_info); - _itr_dim_stat == ROCPROFILER_STATUS_SUCCESS) - { - _dim_ids.reserve(_dim_info.size()); - for(auto ditr : _dim_info) - _dim_ids.emplace_back(ditr.id); - } - else - { - ROCP_WARNING << fmt::format("rocprofiler_iterate_counter_dimensions(...) " - "for counter {} returned {} :: {}", - _info.name, - rocprofiler_get_status_name(_itr_dim_stat), - rocprofiler_get_status_string(_itr_dim_stat)); - } + for(uint64_t j = 0; j < _info.dimensions_count; ++j) + _dim_ids.emplace_back(_info.dimensions[j].id); data_v->at(id).emplace_back( id, _info, std::move(_dim_ids), std::move(_dim_info)); diff --git a/source/lib/rocprofiler-sdk-roctx/abi.cpp b/source/lib/rocprofiler-sdk-roctx/abi.cpp index a3307a4906..3becffd758 100644 --- a/source/lib/rocprofiler-sdk-roctx/abi.cpp +++ b/source/lib/rocprofiler-sdk-roctx/abi.cpp @@ -27,7 +27,7 @@ #include #include #include -#include +#include namespace rocprofiler { diff --git a/source/lib/rocprofiler-sdk-roctx/roctx.cpp b/source/lib/rocprofiler-sdk-roctx/roctx.cpp index 5fb1f61e88..f1bc12cd1d 100644 --- a/source/lib/rocprofiler-sdk-roctx/roctx.cpp +++ b/source/lib/rocprofiler-sdk-roctx/roctx.cpp @@ -143,7 +143,7 @@ get_table_impl() { rocprofiler::common::init_logging("ROCTX"); - auto*& tbl = rocprofiler::common::static_object::construct(); + auto*& tbl = common::static_object::construct(); tbl->core = roctxCoreApiTable_t{sizeof(roctxCoreApiTable_t), &::rocprofiler::roctx::roctxMarkA, diff --git a/source/lib/rocprofiler-sdk-tool/tool.cpp b/source/lib/rocprofiler-sdk-tool/tool.cpp index 0d3ffb66ff..9214406d0f 100644 --- a/source/lib/rocprofiler-sdk-tool/tool.cpp +++ b/source/lib/rocprofiler-sdk-tool/tool.cpp @@ -531,12 +531,11 @@ hip_stream_display_callback(rocprofiler_callback_tracing_record_t record, rocprofiler_user_data_t* user_data, void* data) { - if(tool::get_config().group_by_queue || - record.kind != ROCPROFILER_CALLBACK_TRACING_HIP_STREAM_API) + if(tool::get_config().group_by_queue || record.kind != ROCPROFILER_CALLBACK_TRACING_HIP_STREAM) return; // Extract stream ID from record auto* stream_handle_data = - static_cast(record.payload); + static_cast(record.payload); auto stream_id = stream_handle_data->stream_id; // STREAM_HANDLE_CREATE and DESTROY are no-ops if(record.operation == ROCPROFILER_HIP_STREAM_CREATE) @@ -996,7 +995,7 @@ buffered_tracing_callback(rocprofiler_context_id_t /*context*/, using counter_vec_t = std::vector; using agent_counter_map_t = - std::unordered_map>; + std::unordered_map>; auto get_gpu_agents() @@ -1014,16 +1013,16 @@ struct agent_profiles { std::unordered_map> current_iter; const uint64_t rotation; - const std::unordered_map> + const std::unordered_map> profiles; }; -std::optional +std::optional construct_counter_collection_profile(rocprofiler_agent_id_t agent_id, const std::set& counters) { static const auto gpu_agents_counter_info = get_agent_counter_info(); - auto profile = std::optional{}; + auto profile = std::optional{}; auto counters_v = counter_vec_t{}; auto found_v = std::vector{}; const auto* agent_v = tool_metadata->get_agent(agent_id); @@ -1078,8 +1077,8 @@ construct_counter_collection_profile(rocprofiler_agent_id_t agent_id, if(!counters_v.empty()) { - auto profile_v = rocprofiler_profile_config_id_t{}; - ROCPROFILER_CALL(rocprofiler_create_profile_config( + auto profile_v = rocprofiler_counter_config_id_t{}; + ROCPROFILER_CALL(rocprofiler_create_counter_config( agent_id, counters_v.data(), counters_v.size(), &profile_v), "Could not construct profile cfg"); profile = profile_v; @@ -1090,7 +1089,7 @@ construct_counter_collection_profile(rocprofiler_agent_id_t agent_id, agent_profiles generate_agent_profiles() { - std::unordered_map> + std::unordered_map> profiles; std::unordered_map> pos; for(const auto& agent : get_gpu_agents()) @@ -1110,7 +1109,7 @@ generate_agent_profiles() } // this function creates a rocprofiler profile config on the first entry -std::optional +std::optional get_device_counting_service(rocprofiler_agent_id_t agent_id) { static auto agent_profiles = generate_agent_profiles(); @@ -1292,7 +1291,7 @@ att_dispatch_callback(rocprofiler_agent_id_t /* agent_id */, void dispatch_callback(rocprofiler_dispatch_counting_service_data_t dispatch_data, - rocprofiler_profile_config_id_t* config, + rocprofiler_counter_config_id_t* config, rocprofiler_user_data_t* user_data, void* /*callback_data_args*/) { @@ -1804,14 +1803,14 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data) ROCPROFILER_CALL(rocprofiler_create_context(&hip_stream_display_ctx), "failed to create context"); - ROCPROFILER_CALL(rocprofiler_configure_callback_tracing_service( - hip_stream_display_ctx, - ROCPROFILER_CALLBACK_TRACING_HIP_STREAM_API, - nullptr, - 0, - hip_stream_display_callback, - nullptr), - "stream tracing configure failed"); + ROCPROFILER_CALL( + rocprofiler_configure_callback_tracing_service(hip_stream_display_ctx, + ROCPROFILER_CALLBACK_TRACING_HIP_STREAM, + nullptr, + 0, + hip_stream_display_callback, + nullptr), + "stream tracing configure failed"); ROCPROFILER_CALL(rocprofiler_start_context(hip_stream_display_ctx), "start context failed"); } if(tool::get_config().kernel_rename || !tool::get_config().group_by_queue) diff --git a/source/lib/rocprofiler-sdk/CMakeLists.txt b/source/lib/rocprofiler-sdk/CMakeLists.txt index 134cd47ea5..880645a7c6 100644 --- a/source/lib/rocprofiler-sdk/CMakeLists.txt +++ b/source/lib/rocprofiler-sdk/CMakeLists.txt @@ -10,19 +10,19 @@ set(ROCPROFILER_LIB_SOURCES agent.cpp buffer.cpp buffer_tracing.cpp - device_counting_service.cpp callback_tracing.cpp context.cpp + counter_config.cpp counters.cpp - dispatch_profile.cpp + device_counting_service.cpp + dispatch_counting_service.cpp external_correlation.cpp intercept_table.cpp internal_threading.cpp ompt.cpp pc_sampling.cpp - profile_config.cpp - rocprofiler.cpp registration.cpp + rocprofiler.cpp runtime_initialization.cpp) # ----------------------------------------------------------------------------------------# diff --git a/source/lib/rocprofiler-sdk/buffer_tracing.cpp b/source/lib/rocprofiler-sdk/buffer_tracing.cpp index dbee75341b..fc0606f603 100644 --- a/source/lib/rocprofiler-sdk/buffer_tracing.cpp +++ b/source/lib/rocprofiler-sdk/buffer_tracing.cpp @@ -99,7 +99,7 @@ ROCPROFILER_BUFFER_TRACING_KIND_STRING(OMPT) ROCPROFILER_BUFFER_TRACING_KIND_STRING(RUNTIME_INITIALIZATION) ROCPROFILER_BUFFER_TRACING_KIND_STRING(ROCDECODE_API) ROCPROFILER_BUFFER_TRACING_KIND_STRING(ROCJPEG_API) -ROCPROFILER_BUFFER_TRACING_KIND_STRING(HIP_STREAM_API) +ROCPROFILER_BUFFER_TRACING_KIND_STRING(HIP_STREAM) ROCPROFILER_BUFFER_TRACING_KIND_STRING(HIP_RUNTIME_API_EXT) ROCPROFILER_BUFFER_TRACING_KIND_STRING(HIP_COMPILER_API_EXT) @@ -312,7 +312,7 @@ rocprofiler_query_buffer_tracing_kind_operation_name(rocprofiler_buffer_tracing_ val = rocprofiler::rocjpeg::name_by_id(operation); break; } - case ROCPROFILER_BUFFER_TRACING_HIP_STREAM_API: + case ROCPROFILER_BUFFER_TRACING_HIP_STREAM: { val = rocprofiler::hip::stream::name_by_id(operation); break; @@ -460,7 +460,7 @@ rocprofiler_iterate_buffer_tracing_kind_operations( ops = rocprofiler::rocjpeg::get_ids(); break; } - case ROCPROFILER_BUFFER_TRACING_HIP_STREAM_API: + case ROCPROFILER_BUFFER_TRACING_HIP_STREAM: { ops = rocprofiler::hip::stream::get_ids(); break; diff --git a/source/lib/rocprofiler-sdk/callback_tracing.cpp b/source/lib/rocprofiler-sdk/callback_tracing.cpp index 153055d0d4..e48cc32500 100644 --- a/source/lib/rocprofiler-sdk/callback_tracing.cpp +++ b/source/lib/rocprofiler-sdk/callback_tracing.cpp @@ -95,7 +95,7 @@ ROCPROFILER_CALLBACK_TRACING_KIND_STRING(OMPT) ROCPROFILER_CALLBACK_TRACING_KIND_STRING(RUNTIME_INITIALIZATION) ROCPROFILER_CALLBACK_TRACING_KIND_STRING(ROCDECODE_API) ROCPROFILER_CALLBACK_TRACING_KIND_STRING(ROCJPEG_API) -ROCPROFILER_CALLBACK_TRACING_KIND_STRING(HIP_STREAM_API) +ROCPROFILER_CALLBACK_TRACING_KIND_STRING(HIP_STREAM) template std::pair @@ -290,7 +290,7 @@ rocprofiler_query_callback_tracing_kind_operation_name(rocprofiler_callback_trac val = rocprofiler::rocjpeg::name_by_id(operation); break; } - case ROCPROFILER_CALLBACK_TRACING_HIP_STREAM_API: + case ROCPROFILER_CALLBACK_TRACING_HIP_STREAM: { val = rocprofiler::hip::stream::name_by_id(operation); break; @@ -434,7 +434,7 @@ rocprofiler_iterate_callback_tracing_kind_operations( ops = rocprofiler::rocjpeg::get_ids(); break; } - case ROCPROFILER_CALLBACK_TRACING_HIP_STREAM_API: + case ROCPROFILER_CALLBACK_TRACING_HIP_STREAM: { ops = rocprofiler::hip::stream::get_ids(); break; @@ -583,7 +583,7 @@ rocprofiler_iterate_callback_tracing_kind_operation_args( case ROCPROFILER_CALLBACK_TRACING_RUNTIME_INITIALIZATION: case ROCPROFILER_CALLBACK_TRACING_ROCDECODE_API: case ROCPROFILER_CALLBACK_TRACING_ROCJPEG_API: - case ROCPROFILER_CALLBACK_TRACING_HIP_STREAM_API: + case ROCPROFILER_CALLBACK_TRACING_HIP_STREAM: { return ROCPROFILER_STATUS_ERROR_NOT_IMPLEMENTED; } diff --git a/source/lib/rocprofiler-sdk/profile_config.cpp b/source/lib/rocprofiler-sdk/counter_config.cpp similarity index 91% rename from source/lib/rocprofiler-sdk/profile_config.cpp rename to source/lib/rocprofiler-sdk/counter_config.cpp index 2462efee97..06a446624a 100644 --- a/source/lib/rocprofiler-sdk/profile_config.cpp +++ b/source/lib/rocprofiler-sdk/counter_config.cpp @@ -44,17 +44,17 @@ extern "C" { * @return ::rocprofiler_status_t */ rocprofiler_status_t -rocprofiler_create_profile_config(rocprofiler_agent_id_t agent_id, +rocprofiler_create_counter_config(rocprofiler_agent_id_t agent_id, rocprofiler_counter_id_t* counters_list, size_t counters_count, - rocprofiler_profile_config_id_t* config_id) + rocprofiler_counter_config_id_t* config_id) { std::unordered_set already_added; const auto* agent = ::rocprofiler::agent::get_agent(agent_id); if(!agent) return ROCPROFILER_STATUS_ERROR_AGENT_NOT_FOUND; - std::shared_ptr config = - std::make_shared(); + std::shared_ptr config = + std::make_shared(); auto metrics_map = rocprofiler::counters::loadMetrics(); const auto& id_map = metrics_map->id_to_metric; @@ -78,7 +78,7 @@ rocprofiler_create_profile_config(rocprofiler_agent_id_t agent_id, if(config_id->handle != 0) { // Copy existing counters from previous config - if(auto existing = rocprofiler::counters::get_profile_config(*config_id)) + if(auto existing = rocprofiler::counters::get_counter_config(*config_id)) { for(const auto& metric : existing->metrics) { @@ -100,7 +100,7 @@ rocprofiler_create_profile_config(rocprofiler_agent_id_t agent_id, } rocprofiler_status_t -rocprofiler_destroy_profile_config(rocprofiler_profile_config_id_t config_id) +rocprofiler_destroy_counter_config(rocprofiler_counter_config_id_t config_id) { rocprofiler::counters::destroy_counter_profile(config_id.handle); return ROCPROFILER_STATUS_SUCCESS; diff --git a/source/lib/rocprofiler-sdk/counters.cpp b/source/lib/rocprofiler-sdk/counters.cpp index 9b63ab6f76..409fd5fe74 100644 --- a/source/lib/rocprofiler-sdk/counters.cpp +++ b/source/lib/rocprofiler-sdk/counters.cpp @@ -23,28 +23,61 @@ #include #include #include +#include #include #include "lib/common/container/small_vector.hpp" #include "lib/common/logging.hpp" +#include "lib/common/static_object.hpp" #include "lib/common/string_entry.hpp" -#include "lib/common/utility.hpp" +#include "lib/common/synchronized.hpp" #include "lib/rocprofiler-sdk/agent.hpp" #include "lib/rocprofiler-sdk/counters/dimensions.hpp" #include "lib/rocprofiler-sdk/counters/evaluate_ast.hpp" #include "lib/rocprofiler-sdk/counters/id_decode.hpp" #include "lib/rocprofiler-sdk/counters/metrics.hpp" +#include "lib/rocprofiler-sdk/hsa/agent_cache.hpp" +#include "lib/rocprofiler-sdk/hsa/queue.hpp" +#include "lib/rocprofiler-sdk/hsa/queue_controller.hpp" +namespace rocprofiler +{ +namespace counters +{ namespace { const char* -get_static_string(const std::string& str) +get_static_string(std::string_view str) { - return rocprofiler::common::get_string_entry(rocprofiler::common::add_string_entry(str)) - ->c_str(); + return common::get_string_entry(str)->c_str(); +} + +template +const std::vector* +get_static_ptr(const std::vector& vec) +{ + // The use of std::map is purposeful. Keys can be vectors in map and cannot be in unordered_map. + // Simplifying the code to create these static objects. Given that they are not created often ( + // or looked up often), the performance difference between map and unordered_map is negligible. + using static_ptr_map = std::map, std::unique_ptr>>; + static auto*& static_ptrs = + common::static_object>::construct(); + return static_ptrs->wlock([&](auto& data) { + if(auto it = data.find(vec); it != data.end()) + { + return it->second.get(); + } + data[vec] = std::make_unique>(vec); + return data[vec].get(); + }); } } // namespace +} // namespace counters +} // namespace rocprofiler + +namespace counters = ::rocprofiler::counters; +namespace common = ::rocprofiler::common; extern "C" { /** @@ -64,27 +97,132 @@ rocprofiler_query_counter_info(rocprofiler_counter_id_t counter_id, rocprofiler_counter_info_version_id_t version, void* info) { - if(version != ROCPROFILER_COUNTER_INFO_VERSION_0) - return ROCPROFILER_STATUS_ERROR_INCOMPATIBLE_ABI; - auto metrics_map = rocprofiler::counters::loadMetrics(); + auto metrics_map = counters::loadMetrics(); + const auto& id_map = metrics_map->id_to_metric; - const auto& id_map = metrics_map->id_to_metric; + auto base_info = [&](auto& out_struct) { + if(const auto* metric_ptr = common::get_val(id_map, counter_id.handle)) + { + out_struct.id = counter_id; + out_struct.is_constant = (metric_ptr->constant().empty()) ? 0 : 1; + out_struct.is_derived = (metric_ptr->expression().empty()) ? 0 : 1; + out_struct.name = counters::get_static_string(metric_ptr->name()); + out_struct.description = counters::get_static_string(metric_ptr->description()); + out_struct.block = counters::get_static_string(metric_ptr->block()); + out_struct.expression = counters::get_static_string(metric_ptr->expression()); + return true; + } + return false; + }; - auto& out_struct = *static_cast(info); + auto dim_info = [&](auto& out_struct) { + auto dim_ptr = counters::get_dimension_cache(); - if(const auto* metric_ptr = rocprofiler::common::get_val(id_map, counter_id.handle)) + const auto* dims = common::get_val(dim_ptr->id_to_dim, counter_id.handle); + if(!dims) return false; + + auto _dim_info = std::vector{}; + for(const auto& metric_dim : *dims) + { + _dim_info.emplace_back(rocprofiler_counter_record_dimension_info_t{ + .name = counters::get_static_string(metric_dim.name()), + .instance_size = metric_dim.size(), + .id = static_cast(metric_dim.type())}); + } + + if(_dim_info.empty()) + { + // Can be 0 if the counter is not known by AQLProfile. This is the case + // if it was added in a later version of AQLProfile. + out_struct.dimensions = nullptr; + out_struct.dimensions_count = 0; + return true; + } + + out_struct.dimensions = counters::get_static_ptr(_dim_info)->data(); + out_struct.dimensions_count = _dim_info.size(); + return true; + }; + + // Construct all possible permutations of instance ids. This is every instance + // that can be returned by the counter across all dimensions. + auto dim_permutations = [&](auto& out_struct) { + auto dim_ptr = counters::get_dimension_cache(); + + const auto* dims = common::get_val(dim_ptr->id_to_dim, counter_id.handle); + if(!dims) return false; + + std::vector instances; + + for(const auto& metric_dim : *dims) + { + if(metric_dim.size() == 0) continue; + std::vector tmp; + // If no instances are found, create the first set of instances + if(instances.empty()) + { + for(size_t i = 0; i < metric_dim.size(); i++) + { + auto& rec = instances.emplace_back(); + counters::set_dim_in_rec(rec, metric_dim.type(), i); + counters::set_counter_in_rec(rec, counter_id); + } + } + else + { + // For each instance, create a new set of instances with the new dimension added. + // This will create all possible permutations of the dimensions. + for(size_t i = 0; i < metric_dim.size(); i++) + { + for(const auto& instance : instances) + { + auto& rec = tmp.emplace_back(instance); + counters::set_dim_in_rec(rec, metric_dim.type(), i); + counters::set_counter_in_rec(rec, counter_id); + } + } + instances = tmp; + } + } + if(instances.empty()) + { + out_struct.instance_ids = nullptr; + out_struct.instance_ids_count = 0; + return true; + } + + out_struct.instance_ids = counters::get_static_ptr(instances)->data(); + out_struct.instance_ids_count = instances.size(); + return true; + }; + + switch(version) { - out_struct.id = counter_id; - out_struct.is_constant = (metric_ptr->constant().empty()) ? 0 : 1; - out_struct.is_derived = (metric_ptr->expression().empty()) ? 0 : 1; - out_struct.name = get_static_string(metric_ptr->name()); - out_struct.description = get_static_string(metric_ptr->description()); - out_struct.block = get_static_string(metric_ptr->block()); - out_struct.expression = get_static_string(metric_ptr->expression()); - return ROCPROFILER_STATUS_SUCCESS; + case ROCPROFILER_COUNTER_INFO_VERSION_0: + { + auto& _out_struct = *static_cast(info); + + if(base_info(_out_struct)) return ROCPROFILER_STATUS_SUCCESS; + return ROCPROFILER_STATUS_ERROR_COUNTER_NOT_FOUND; + } + break; + case ROCPROFILER_COUNTER_INFO_VERSION_1: + { + auto& _out_struct = *static_cast(info); + + if(!base_info(_out_struct)) return ROCPROFILER_STATUS_ERROR_COUNTER_NOT_FOUND; + if(!dim_info(_out_struct)) return ROCPROFILER_STATUS_ERROR_DIM_NOT_FOUND; + if(!dim_permutations(_out_struct)) return ROCPROFILER_STATUS_ERROR_DIM_NOT_FOUND; + + return ROCPROFILER_STATUS_SUCCESS; + } + break; + default: + { + return ROCPROFILER_STATUS_ERROR_INCOMPATIBLE_ABI; + } } - ROCP_ERROR << fmt::format("Could not find counter with id = {}", counter_id.handle); return ROCPROFILER_STATUS_ERROR_COUNTER_NOT_FOUND; } @@ -102,9 +240,9 @@ rocprofiler_query_counter_instance_count(rocprofiler_agent_id_t, size_t* instance_count) { *instance_count = 0; - auto dim_ptr = rocprofiler::counters::get_dimension_cache(); + auto dim_ptr = counters::get_dimension_cache(); - const auto* dims = rocprofiler::common::get_val(dim_ptr->id_to_dim, counter_id.handle); + const auto* dims = common::get_val(dim_ptr->id_to_dim, counter_id.handle); if(!dims) return ROCPROFILER_STATUS_ERROR_COUNTER_NOT_FOUND; for(const auto& metric_dim : *dims) @@ -133,7 +271,7 @@ rocprofiler_iterate_agent_supported_counters(rocprofiler_agent_id_t const auto* agent = rocprofiler::agent::get_agent(agent_id); if(!agent) return ROCPROFILER_STATUS_ERROR_AGENT_NOT_FOUND; - auto metrics = rocprofiler::counters::getMetricsForAgent(agent->name); + auto metrics = counters::getMetricsForAgent(agent->name); if(metrics.empty()) return ROCPROFILER_STATUS_ERROR_AGENT_ARCH_NOT_SUPPORTED; std::vector ids; @@ -158,7 +296,7 @@ rocprofiler_query_record_counter_id(rocprofiler_counter_instance_id_t id, rocprofiler_counter_id_t* counter_id) { // Get counter id from record - *counter_id = rocprofiler::counters::rec_to_counter_id(id); + *counter_id = counters::rec_to_counter_id(id); return ROCPROFILER_STATUS_SUCCESS; } @@ -167,8 +305,8 @@ rocprofiler_query_record_dimension_position(rocprofiler_counter_instance_id_t i rocprofiler_counter_dimension_id_t dim, size_t* pos) { - *pos = rocprofiler::counters::rec_to_dim_pos( - id, static_cast(dim)); + *pos = counters::rec_to_dim_pos( + id, static_cast(dim)); return ROCPROFILER_STATUS_SUCCESS; } @@ -177,17 +315,17 @@ rocprofiler_iterate_counter_dimensions(rocprofiler_counter_id_t id, rocprofiler_available_dimensions_cb_t info_cb, void* user_data) { - auto dim_ptr = rocprofiler::counters::get_dimension_cache(); + auto dim_ptr = counters::get_dimension_cache(); - const auto* dims = rocprofiler::common::get_val(dim_ptr->id_to_dim, id.handle); + const auto* dims = common::get_val(dim_ptr->id_to_dim, id.handle); if(!dims) return ROCPROFILER_STATUS_ERROR_COUNTER_NOT_FOUND; // This is likely faster than a map lookup given the limited number of dims. - rocprofiler::common::container::small_vector user_dims; + auto user_dims = common::container::small_vector{}; for(const auto& internal_dim : *dims) { auto& dim = user_dims.emplace_back(); - dim.name = get_static_string(internal_dim.name()); + dim.name = counters::get_static_string(internal_dim.name()); dim.instance_size = internal_dim.size(); dim.id = static_cast(internal_dim.type()); } @@ -205,12 +343,12 @@ rocprofiler_iterate_counter_dimensions(rocprofiler_counter_id_t id, rocprofiler_status_t rocprofiler_load_counter_definition(const char* yaml, size_t size, rocprofiler_counter_flag_t flags) { - rocprofiler::counters::CustomCounterDefinition def; + counters::CustomCounterDefinition def; if(yaml == nullptr && size != 0) return ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT; def.data = std::string(yaml, size); def.append = (flags == ROCPROFILER_COUNTER_FLAG_APPEND_DEFINITION ? true : false); def.loaded = false; - return rocprofiler::counters::setCustomCounterDefinition(def); + return counters::setCustomCounterDefinition(def); } rocprofiler_status_t @@ -226,25 +364,23 @@ rocprofiler_create_counter(const char* name, const auto* agent_ptr = rocprofiler::agent::get_agent(agent); if(!agent_ptr) return ROCPROFILER_STATUS_ERROR_AGENT_NOT_FOUND; - rocprofiler::counters::Metric new_metric( - "", - std::string(name, name_len), - "", - "", - std::string((description ? description : ""), description_len), - std::string(expr, expr_len), - "", - -1); + counters::Metric new_metric("", + std::string(name, name_len), + "", + "", + std::string((description ? description : ""), description_len), + std::string(expr, expr_len), + "", + -1); // Validate the metric. Checks for duplicate names and invalid expressions. - if(auto status = rocprofiler::counters::check_ast_generation(agent_ptr->name, new_metric); + if(auto status = counters::check_ast_generation(agent_ptr->name, new_metric); status != ROCPROFILER_STATUS_SUCCESS) { return status; } - auto add_metric = - rocprofiler::counters::loadMetrics(true, std::make_pair(agent_ptr->name, new_metric)); + auto add_metric = counters::loadMetrics(true, std::make_pair(agent_ptr->name, new_metric)); if(add_metric->arch_to_metric.at(agent_ptr->name).back().name() != new_metric.name()) { @@ -256,8 +392,8 @@ rocprofiler_create_counter(const char* name, // Regenerate ASTs and Dimension Cache try { - rocprofiler::counters::get_ast_map(true); - rocprofiler::counters::get_dimension_cache(true); + counters::get_ast_map(true); + counters::get_dimension_cache(true); } catch(std::exception& e) { ROCP_FATAL << "Could not regenerate ASTs and Dimension Cache " << e.what(); diff --git a/source/lib/rocprofiler-sdk/counters/controller.cpp b/source/lib/rocprofiler-sdk/counters/controller.cpp index 30445cfb93..2a81cfdd4c 100644 --- a/source/lib/rocprofiler-sdk/counters/controller.cpp +++ b/source/lib/rocprofiler-sdk/counters/controller.cpp @@ -46,12 +46,12 @@ CounterController::CounterController() // Note: these profiles can be used across multiple contexts // and are independent of the context. uint64_t -CounterController::add_profile(std::shared_ptr&& config) +CounterController::add_profile(std::shared_ptr&& config) { static std::atomic profile_val = 1; uint64_t ret = 0; _configs.wlock([&](auto& data) { - config->id = rocprofiler_profile_config_id_t{.handle = profile_val}; + config->id = rocprofiler_counter_config_id_t{.handle = profile_val}; data.emplace(profile_val, std::move(config)); ret = profile_val; profile_val++; @@ -66,11 +66,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_device_counting_service_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_cb_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; @@ -127,13 +127,12 @@ CounterController::configure_agent_collection(rocprofiler_context_id_t context_i // the AQL packet generator for injecting packets. Note: the service is created // in the stop state. rocprofiler_status_t -CounterController::configure_dispatch( - rocprofiler_context_id_t context_id, - rocprofiler_buffer_id_t buffer, - rocprofiler_dispatch_counting_service_callback_t callback, - void* callback_args, - rocprofiler_profile_counting_record_callback_t record_callback, - void* record_callback_args) +CounterController::configure_dispatch(rocprofiler_context_id_t context_id, + rocprofiler_buffer_id_t buffer, + rocprofiler_dispatch_counting_service_cb_t callback, + void* callback_args, + rocprofiler_dispatch_counting_record_cb_t record_callback, + void* record_callback_args) { auto* ctx_p = rocprofiler::context::get_mutable_registered_context(context_id); if(!ctx_p) return ROCPROFILER_STATUS_ERROR_CONTEXT_INVALID; @@ -169,10 +168,10 @@ CounterController::configure_dispatch( return ROCPROFILER_STATUS_SUCCESS; } -std::shared_ptr -CounterController::get_profile_cfg(rocprofiler_profile_config_id_t id) +std::shared_ptr +CounterController::get_profile_cfg(rocprofiler_counter_config_id_t id) { - std::shared_ptr cfg; + std::shared_ptr cfg; _configs.rlock([&](const auto& map) { cfg = map.at(id.handle); }); return cfg; } @@ -185,10 +184,10 @@ get_controller() } rocprofiler_status_t -create_counter_profile(std::shared_ptr config) +create_counter_profile(std::shared_ptr config) { auto status = ROCPROFILER_STATUS_SUCCESS; - if(status = counters::counter_callback_info::setup_profile_config(config); + if(status = counters::counter_callback_info::setup_counter_config(config); status != ROCPROFILER_STATUS_SUCCESS) { return status; @@ -210,8 +209,8 @@ destroy_counter_profile(uint64_t id) get_controller().destroy_profile(id); } -std::shared_ptr -get_profile_config(rocprofiler_profile_config_id_t id) +std::shared_ptr +get_counter_config(rocprofiler_counter_config_id_t id) { try { diff --git a/source/lib/rocprofiler-sdk/counters/controller.hpp b/source/lib/rocprofiler-sdk/counters/controller.hpp index 3a0b3a277c..a73b0dc4c1 100644 --- a/source/lib/rocprofiler-sdk/counters/controller.hpp +++ b/source/lib/rocprofiler-sdk/counters/controller.hpp @@ -42,7 +42,7 @@ namespace counters // to collect counters on, the metrics to collect, the hw // counters needed to evaluate the metrics, and the ASTs. // This profile can be shared among many rocprof contexts. -struct profile_config +struct counter_config { const rocprofiler_agent_t* agent = nullptr; std::vector metrics{}; @@ -55,13 +55,12 @@ struct profile_config std::set required_special_counters{}; // ASTs to evaluate std::vector asts{}; - rocprofiler_profile_config_id_t id{.handle = 0}; + rocprofiler_counter_config_id_t id{.handle = 0}; // Packet generator to create AQL packets for insertion std::unique_ptr pkt_generator{nullptr}; // A packet cache of AQL packets. This allows reuse of AQL packets (preventing costly // allocation of new packets/destruction). - rocprofiler::common::Synchronized>> - packets{}; + common::Synchronized>> packets{}; }; class CounterController @@ -72,7 +71,7 @@ public: // Adds a counter collection profile to our global cache. // Note: these profiles can be used across multiple contexts // and are independent of the context. - uint64_t add_profile(std::shared_ptr&& config); + uint64_t add_profile(std::shared_ptr&& config); void destroy_profile(uint64_t id); // Setup the counter collection service. counter_callback_info is created here @@ -80,37 +79,36 @@ public: // the AQL packet generator for injecting packets. Note: the service is created // in the stop state. static rocprofiler_status_t configure_dispatch( - rocprofiler_context_id_t context_id, - rocprofiler_buffer_id_t buffer, - 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); + rocprofiler_context_id_t context_id, + rocprofiler_buffer_id_t buffer, + rocprofiler_dispatch_counting_service_cb_t callback, + void* callback_args, + rocprofiler_dispatch_counting_record_cb_t record_callback, + void* record_callback_args); + std::shared_ptr get_profile_cfg(rocprofiler_counter_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_device_counting_service_callback_t cb, - void* user_data); + rocprofiler_context_id_t context_id, + rocprofiler_buffer_id_t buffer_id, + rocprofiler_agent_id_t agent_id, + rocprofiler_device_counting_service_cb_t cb, + void* user_data); private: - rocprofiler::common::Synchronized>> - _configs; + common::Synchronized>> _configs; }; CounterController& get_controller(); rocprofiler_status_t -create_counter_profile(std::shared_ptr config); +create_counter_profile(std::shared_ptr config); void destroy_counter_profile(uint64_t id); -std::shared_ptr -get_profile_config(rocprofiler_profile_config_id_t id); +std::shared_ptr +get_counter_config(rocprofiler_counter_config_id_t id); } // namespace counters } // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk/counters/core.cpp b/source/lib/rocprofiler-sdk/counters/core.cpp index 5bb00cebb1..1f3928659d 100644 --- a/source/lib/rocprofiler-sdk/counters/core.cpp +++ b/source/lib/rocprofiler-sdk/counters/core.cpp @@ -40,7 +40,7 @@ namespace rocprofiler namespace counters { rocprofiler_status_t -counter_callback_info::setup_profile_config(std::shared_ptr& profile) +counter_callback_info::setup_counter_config(std::shared_ptr& profile) { if(profile->pkt_generator || !profile->reqired_hw_counters.empty()) { @@ -115,12 +115,12 @@ counter_callback_info::setup_profile_config(std::shared_ptr& pro rocprofiler_status_t counter_callback_info::get_packet(std::unique_ptr& ret_pkt, - std::shared_ptr& profile) + std::shared_ptr& profile) { rocprofiler_status_t status; // Check packet cache profile->packets.wlock([&](auto& pkt_vector) { - status = counter_callback_info::setup_profile_config(profile); + status = counter_callback_info::setup_counter_config(profile); if(!pkt_vector.empty() && status == ROCPROFILER_STATUS_SUCCESS) { ret_pkt = std::move(pkt_vector.back()); @@ -216,21 +216,21 @@ 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_device_counting_service_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_cb_t cb, + void* user_data) { return get_controller().configure_agent_collection( context_id, buffer_id, agent_id, cb, user_data); } rocprofiler_status_t -configure_buffered_dispatch(rocprofiler_context_id_t context_id, - rocprofiler_buffer_id_t buffer, - rocprofiler_dispatch_counting_service_callback_t callback, - void* callback_args) +configure_buffered_dispatch(rocprofiler_context_id_t context_id, + rocprofiler_buffer_id_t buffer, + rocprofiler_dispatch_counting_service_cb_t callback, + void* callback_args) { CHECK_NE(buffer.handle, 0); return get_controller().configure_dispatch( @@ -238,11 +238,11 @@ configure_buffered_dispatch(rocprofiler_context_id_t con } rocprofiler_status_t -configure_callback_dispatch(rocprofiler_context_id_t context_id, - rocprofiler_dispatch_counting_service_callback_t callback, - void* callback_data_args, - rocprofiler_profile_counting_record_callback_t record_callback, - void* record_callback_args) +configure_callback_dispatch(rocprofiler_context_id_t context_id, + rocprofiler_dispatch_counting_service_cb_t callback, + void* callback_data_args, + rocprofiler_dispatch_counting_record_cb_t record_callback, + void* record_callback_args) { return get_controller().configure_dispatch(context_id, {.handle = 0}, diff --git a/source/lib/rocprofiler-sdk/counters/core.hpp b/source/lib/rocprofiler-sdk/counters/core.hpp index 8017fd5ba2..588098ceda 100644 --- a/source/lib/rocprofiler-sdk/counters/core.hpp +++ b/source/lib/rocprofiler-sdk/counters/core.hpp @@ -45,7 +45,7 @@ namespace counters struct counter_callback_info { // User callback - rocprofiler_dispatch_counting_service_callback_t user_cb{nullptr}; + rocprofiler_dispatch_counting_service_cb_t user_cb{nullptr}; // User id void* callback_args{nullptr}; // Link to the context this is associated with @@ -58,45 +58,45 @@ struct counter_callback_info // Buffer to use for storing counter data. Used if callback is not set. std::optional buffer; - rocprofiler_profile_counting_record_callback_t record_callback; - void* record_callback_args; + rocprofiler_dispatch_counting_record_cb_t record_callback; + void* record_callback_args; // Facilitates the return of an AQL Packet to the profile config that constructed it. - rocprofiler::common::Synchronized< - std::unordered_map>> + common::Synchronized< + std::unordered_map>> packet_return_map{}; - static rocprofiler_status_t setup_profile_config(std::shared_ptr&); + static rocprofiler_status_t setup_counter_config(std::shared_ptr&); rocprofiler_status_t get_packet(std::unique_ptr&, - std::shared_ptr&); + std::shared_ptr&); }; uint64_t -create_counter_profile(std::shared_ptr&& config); +create_counter_profile(std::shared_ptr&& config); void destroy_counter_profile(uint64_t id); rocprofiler_status_t -configure_buffered_dispatch(rocprofiler_context_id_t context_id, - rocprofiler_buffer_id_t buffer, - rocprofiler_dispatch_counting_service_callback_t callback, - void* callback_args); +configure_buffered_dispatch(rocprofiler_context_id_t context_id, + rocprofiler_buffer_id_t buffer, + rocprofiler_dispatch_counting_service_cb_t callback, + void* callback_args); rocprofiler_status_t -configure_callback_dispatch(rocprofiler_context_id_t context_id, - rocprofiler_dispatch_counting_service_callback_t callback, - void* callback_data_args, - rocprofiler_profile_counting_record_callback_t record_callback, - void* record_callback_args); +configure_callback_dispatch(rocprofiler_context_id_t context_id, + rocprofiler_dispatch_counting_service_cb_t callback, + void* callback_data_args, + rocprofiler_dispatch_counting_record_cb_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_device_counting_service_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_cb_t cb, + void* user_data); void start_context(const context::context*); diff --git a/source/lib/rocprofiler-sdk/counters/device_counting.cpp b/source/lib/rocprofiler-sdk/counters/device_counting.cpp index 3b36657868..91cb0c2c87 100644 --- a/source/lib/rocprofiler-sdk/counters/device_counting.cpp +++ b/source/lib/rocprofiler-sdk/counters/device_counting.cpp @@ -103,9 +103,9 @@ header_pkt(hsa_packet_type_t type) } std::unique_ptr -construct_aql_pkt(std::shared_ptr& profile) +construct_aql_pkt(std::shared_ptr& profile) { - if(counter_callback_info::setup_profile_config(profile) != ROCPROFILER_STATUS_SUCCESS) + if(counter_callback_info::setup_counter_config(profile) != ROCPROFILER_STATUS_SUCCESS) { return nullptr; } @@ -422,11 +422,11 @@ start_agent_ctx(const context::context* ctx) {.handle = ctx->context_idx}, callback_data.agent_id, [](rocprofiler_context_id_t context_id, - rocprofiler_profile_config_id_t config_id) -> rocprofiler_status_t { + rocprofiler_counter_config_id_t config_id) -> rocprofiler_status_t { auto* cb_ctx = rocprofiler::context::get_mutable_registered_context(context_id); if(!cb_ctx) return ROCPROFILER_STATUS_ERROR_CONTEXT_INVALID; - auto config = rocprofiler::counters::get_profile_config(config_id); + auto config = rocprofiler::counters::get_counter_config(config_id); if(!config) return ROCPROFILER_STATUS_ERROR_PROFILE_NOT_FOUND; if(!cb_ctx->device_counter_collection) diff --git a/source/lib/rocprofiler-sdk/counters/device_counting.hpp b/source/lib/rocprofiler-sdk/counters/device_counting.hpp index 52ef41e04d..9e8f61aa11 100644 --- a/source/lib/rocprofiler-sdk/counters/device_counting.hpp +++ b/source/lib/rocprofiler-sdk/counters/device_counting.hpp @@ -38,7 +38,7 @@ struct context; namespace counters { -struct profile_config; +struct counter_config; struct agent_callback_data { @@ -56,9 +56,9 @@ struct agent_callback_data rocprofiler_user_data_t user_data = {.value = 0}; rocprofiler_user_data_t callback_data = {.value = 0}; - std::shared_ptr profile = {}; + std::shared_ptr profile = {}; rocprofiler_agent_id_t agent_id = {.handle = 0}; - rocprofiler_device_counting_service_callback_t cb = nullptr; + rocprofiler_device_counting_service_cb_t cb = nullptr; rocprofiler_buffer_id_t buffer = {.handle = 0}; bool set_profile = false; std::vector* cached_counters = nullptr; diff --git a/source/lib/rocprofiler-sdk/counters/dispatch_handlers.cpp b/source/lib/rocprofiler-sdk/counters/dispatch_handlers.cpp index 17e09d0f8d..1b7cb3301c 100644 --- a/source/lib/rocprofiler-sdk/counters/dispatch_handlers.cpp +++ b/source/lib/rocprofiler-sdk/counters/dispatch_handlers.cpp @@ -91,7 +91,7 @@ queue_cb(const context::context* ctx, } } - auto req_profile = rocprofiler_profile_config_id_t{.handle = 0}; + auto req_profile = rocprofiler_counter_config_id_t{.handle = 0}; auto dispatch_data = common::init_public_api_struct(rocprofiler_dispatch_counting_service_data_t{}); @@ -147,7 +147,7 @@ completed_cb(const context::context* ctx, { CHECK(info && ctx); - std::shared_ptr prof_config; + std::shared_ptr prof_config; // Get the Profile Config std::unique_ptr pkt = nullptr; info->packet_return_map.wlock([&](auto& data) { diff --git a/source/lib/rocprofiler-sdk/counters/sample_processing.hpp b/source/lib/rocprofiler-sdk/counters/sample_processing.hpp index 9a6a70f394..4aff58930d 100644 --- a/source/lib/rocprofiler-sdk/counters/sample_processing.hpp +++ b/source/lib/rocprofiler-sdk/counters/sample_processing.hpp @@ -35,7 +35,7 @@ struct completed_cb_params_t std::shared_ptr info; std::shared_ptr session; kernel_dispatch::profiling_time dispatch_time; - std::shared_ptr prof_config; + std::shared_ptr prof_config; std::unique_ptr pkt; }; diff --git a/source/lib/rocprofiler-sdk/counters/tests/core.cpp b/source/lib/rocprofiler-sdk/counters/tests/core.cpp index 7a289ca6c0..fe2442ce66 100644 --- a/source/lib/rocprofiler-sdk/counters/tests/core.cpp +++ b/source/lib/rocprofiler-sdk/counters/tests/core.cpp @@ -168,7 +168,7 @@ buffered_callback(rocprofiler_context_id_t, void null_dispatch_callback(rocprofiler_dispatch_counting_service_data_t, - rocprofiler_profile_config_id_t*, + rocprofiler_counter_config_id_t*, rocprofiler_user_data_t*, void*) {} @@ -209,15 +209,15 @@ TEST(core, check_packet_generation) /** * Check profile construction */ - rocprofiler_profile_config_id_t cfg_id = {.handle = 0}; + rocprofiler_counter_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( - rocprofiler_create_profile_config(agent.get_rocp_agent()->id, &id, 1, &cfg_id), + rocprofiler_create_counter_config(agent.get_rocp_agent()->id, &id, 1, &cfg_id), "Unable to create profile"); - auto profile = counters::get_profile_config(cfg_id); + auto profile = counters::get_counter_config(cfg_id); ASSERT_TRUE(profile); - EXPECT_EQ(counters::counter_callback_info::setup_profile_config(profile), + EXPECT_EQ(counters::counter_callback_info::setup_counter_config(profile), ROCPROFILER_STATUS_SUCCESS) << fmt::format("Could not build profile for {}", metric.name()); @@ -304,7 +304,7 @@ namespace struct expected_dispatch { // To pass back - rocprofiler_profile_config_id_t id = {.handle = 0}; + rocprofiler_counter_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; @@ -312,12 +312,12 @@ struct expected_dispatch rocprofiler_async_correlation_id_t correlation_id = {.internal = 0, .external = {.value = 0}}; rocprofiler_dim3_t workgroup_size = {0, 0, 0}; rocprofiler_dim3_t grid_size = {0, 0, 0}; - rocprofiler_profile_config_id_t* config = nullptr; + rocprofiler_counter_config_id_t* config = nullptr; }; void user_dispatch_cb(rocprofiler_dispatch_counting_service_data_t dispatch_data, - rocprofiler_profile_config_id_t* config, + rocprofiler_counter_config_id_t* config, rocprofiler_user_data_t* user_data, void* callback_data_args) { @@ -402,9 +402,9 @@ TEST(core, check_callbacks) expected_dispatch expected = {}; rocprofiler_counter_id_t id = {.handle = metric.id()}; ROCPROFILER_CALL( - rocprofiler_create_profile_config(agent.get_rocp_agent()->id, &id, 1, &expected.id), + rocprofiler_create_counter_config(agent.get_rocp_agent()->id, &id, 1, &expected.id), "Unable to create profile"); - auto profile = counters::get_profile_config(expected.id); + auto profile = counters::get_counter_config(expected.id); ASSERT_TRUE(profile); std::shared_ptr cb_info = @@ -502,14 +502,14 @@ TEST(core, destroy_counter_profile) expected_dispatch expected = {}; rocprofiler_counter_id_t id = {.handle = metric.id()}; ROCPROFILER_CALL( - rocprofiler_create_profile_config(agent.get_rocp_agent()->id, &id, 1, &expected.id), + rocprofiler_create_counter_config(agent.get_rocp_agent()->id, &id, 1, &expected.id), "Unable to create profile"); - ROCPROFILER_CALL(rocprofiler_destroy_profile_config(expected.id), + ROCPROFILER_CALL(rocprofiler_destroy_counter_config(expected.id), "Could not delete profile id"); /** * Check the profile was actually destroyed */ - auto profile = counters::get_profile_config(expected.id); + auto profile = counters::get_counter_config(expected.id); EXPECT_FALSE(profile); } } @@ -538,7 +538,7 @@ TEST(core, start_stop_buffered_ctx) &opt_buff_id), "Could not create buffer"); - ROCPROFILER_CALL(rocprofiler_configure_buffered_dispatch_counting_service( + ROCPROFILER_CALL(rocprofiler_configure_buffer_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"); @@ -671,16 +671,16 @@ TEST(core, test_profile_incremental) } } - rocprofiler_profile_config_id_t cfg_id = {}; + rocprofiler_counter_config_id_t cfg_id = {}; // Add one counter from each block to incrementally to make sure we can // add them incrementally for(const auto& [block_name, block_metrics] : metric_blocks) { - rocprofiler_profile_config_id_t old_id = cfg_id; + rocprofiler_counter_config_id_t old_id = cfg_id; rocprofiler_counter_id_t id = {.handle = block_metrics.front().id()}; ROCPROFILER_CALL( - rocprofiler_create_profile_config(agent.get_rocp_agent()->id, &id, 1, &cfg_id), + rocprofiler_create_counter_config(agent.get_rocp_agent()->id, &id, 1, &cfg_id), "Unable to create profile incrementally when we should be able to"); EXPECT_NE(old_id.handle, cfg_id.handle) << "We expect that the handle changes this is due to the existing profile being " @@ -697,7 +697,7 @@ TEST(core, test_profile_incremental) */ rocprofiler_counter_id_t id = {.handle = metric.id()}; if(status = - rocprofiler_create_profile_config(agent.get_rocp_agent()->id, &id, 1, &cfg_id); + rocprofiler_create_counter_config(agent.get_rocp_agent()->id, &id, 1, &cfg_id); status != ROCPROFILER_STATUS_SUCCESS) { break; diff --git a/source/lib/rocprofiler-sdk/counters/tests/device_counting.cpp b/source/lib/rocprofiler-sdk/counters/tests/device_counting.cpp index 588ca4e277..9587b219f3 100644 --- a/source/lib/rocprofiler-sdk/counters/tests/device_counting.cpp +++ b/source/lib/rocprofiler-sdk/counters/tests/device_counting.cpp @@ -349,10 +349,10 @@ protected: /** * Check profile construction */ - rocprofiler_profile_config_id_t cfg_id = {.handle = 0}; + rocprofiler_counter_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), + rocprofiler_create_counter_config(agent.get_rocp_agent()->id, &id, 1, &cfg_id), "Unable to create profile"); ROCPROFILER_CALL( @@ -362,12 +362,12 @@ protected: agent.get_rocp_agent()->id, [](rocprofiler_context_id_t context_id, rocprofiler_agent_id_t, - rocprofiler_agent_set_profile_callback_t set_config, - void* user_data) { + rocprofiler_device_counting_agent_cb_t set_config, + void* user_data) { CHECK(user_data); if(auto status = set_config( context_id, - *static_cast(user_data)); + *static_cast(user_data)); status != ROCPROFILER_STATUS_SUCCESS) { ROCP_FATAL << rocprofiler_get_status_string(status); diff --git a/source/lib/rocprofiler-sdk/counters/tests/dimension.cpp b/source/lib/rocprofiler-sdk/counters/tests/dimension.cpp index dd14d1e4cb..5880f10b2f 100644 --- a/source/lib/rocprofiler-sdk/counters/tests/dimension.cpp +++ b/source/lib/rocprofiler-sdk/counters/tests/dimension.cpp @@ -20,6 +20,10 @@ // OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE // SOFTWARE. +// Allow testing of deprecated calls +#pragma GCC diagnostic push +#pragma GCC diagnostic ignored "-Wdeprecated-declarations" + #include "lib/common/utility.hpp" #include "lib/rocprofiler-sdk/agent.hpp" #include "lib/rocprofiler-sdk/aql/packet_construct.hpp" @@ -320,3 +324,4 @@ TEST(dimension, block_dim_test) hsa_shut_down(); } +#pragma GCC diagnostic pop diff --git a/source/lib/rocprofiler-sdk/counters/tests/metrics_test.cpp b/source/lib/rocprofiler-sdk/counters/tests/metrics_test.cpp index fe31afbf65..fa3ea0fc5b 100644 --- a/source/lib/rocprofiler-sdk/counters/tests/metrics_test.cpp +++ b/source/lib/rocprofiler-sdk/counters/tests/metrics_test.cpp @@ -25,12 +25,17 @@ #include #include +#include +#include #include #include "lib/common/logging.hpp" +#include "lib/common/utility.hpp" #include "lib/rocprofiler-sdk/agent.hpp" +#include "lib/rocprofiler-sdk/counters/dimensions.hpp" #include "lib/rocprofiler-sdk/counters/metrics.hpp" +#include "rocprofiler-sdk/fwd.h" namespace { @@ -216,16 +221,57 @@ TEST(metrics, check_public_api_query) const auto& id_map = metrics_map->id_to_metric; for(const auto& [id, metric] : id_map) { - rocprofiler_counter_info_v0_t version; + rocprofiler_counter_info_v1_t info; - ASSERT_EQ( - rocprofiler_query_counter_info( - {.handle = id}, ROCPROFILER_COUNTER_INFO_VERSION_0, static_cast(&version)), - ROCPROFILER_STATUS_SUCCESS); - EXPECT_EQ(std::string(version.name), metric.name()); - EXPECT_EQ(std::string(version.block), metric.block()); - EXPECT_EQ(std::string(version.expression), metric.expression()); - EXPECT_EQ(version.is_derived, !metric.expression().empty()); - EXPECT_EQ(std::string(version.description), metric.description()); + auto dim_ptr = rocprofiler::counters::get_dimension_cache(); + + const auto* dims = rocprofiler::common::get_val(dim_ptr->id_to_dim, metric.id()); + ASSERT_TRUE(dims); + + auto status = rocprofiler_query_counter_info( + {.handle = id}, ROCPROFILER_COUNTER_INFO_VERSION_1, static_cast(&info)); + ASSERT_EQ(status, ROCPROFILER_STATUS_SUCCESS); + EXPECT_EQ(std::string(info.name ? info.name : ""), metric.name()); + EXPECT_EQ(std::string(info.block ? info.block : ""), metric.block()); + EXPECT_EQ(std::string(info.expression ? info.expression : ""), metric.expression()); + EXPECT_EQ(info.is_derived, !metric.expression().empty()); + EXPECT_EQ(std::string(info.description ? info.description : ""), metric.description()); + + EXPECT_EQ(info.dimensions_count, dims->size()); + for(size_t i = 0; i < info.dimensions_count; i++) + { + const auto& dim = dims->at(i); + EXPECT_EQ(dim.size(), info.dimensions[i].instance_size); + EXPECT_EQ(dim.type(), info.dimensions[i].id); + EXPECT_EQ(std::string(info.dimensions[i].name), dim.name()); + } + + size_t instance_count = 0; + // Checks the equality with the old rocprofiler_query_counter_instance_count + for(const auto& metric_dim : *dims) + { + if(instance_count == 0) + instance_count = metric_dim.size(); + else if(metric_dim.size() > 0) + instance_count = metric_dim.size() * instance_count; + } + + EXPECT_EQ(info.instance_ids_count, instance_count); + std::set> dim_permutations; + + for(size_t i = 0; i < info.instance_ids_count; i++) + { + std::vector dim_ids; + ASSERT_EQ(rocprofiler::counters::rec_to_counter_id(info.instance_ids[i]).handle, + metric.id()); + for(const auto& metric_dim : *dims) + { + dim_ids.push_back( + rocprofiler::counters::rec_to_dim_pos(info.instance_ids[i], metric_dim.type())); + } + // Ensure that the premutation is unique + ASSERT_EQ(dim_permutations.insert(dim_ids).second, true); + } + ASSERT_EQ(instance_count, dim_permutations.size()); } } diff --git a/source/lib/rocprofiler-sdk/device_counting_service.cpp b/source/lib/rocprofiler-sdk/device_counting_service.cpp index 89e4d680df..8d1897517d 100644 --- a/source/lib/rocprofiler-sdk/device_counting_service.cpp +++ b/source/lib/rocprofiler-sdk/device_counting_service.cpp @@ -37,11 +37,11 @@ constexpr auto rocprofiler_context_none = ROCPROFILER_CONTEXT_NONE; extern "C" { rocprofiler_status_t -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_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_cb_t cb, + void* user_data) { return rocprofiler::counters::configure_agent_collection( context_id, buffer_id, agent_id, cb, user_data); diff --git a/source/lib/rocprofiler-sdk/dispatch_profile.cpp b/source/lib/rocprofiler-sdk/dispatch_counting_service.cpp similarity index 81% rename from source/lib/rocprofiler-sdk/dispatch_profile.cpp rename to source/lib/rocprofiler-sdk/dispatch_counting_service.cpp index 8231669115..5279513f81 100644 --- a/source/lib/rocprofiler-sdk/dispatch_profile.cpp +++ b/source/lib/rocprofiler-sdk/dispatch_counting_service.cpp @@ -39,11 +39,11 @@ extern "C" { * @return ::rocprofiler_status_t */ rocprofiler_status_t -rocprofiler_configure_buffered_dispatch_counting_service( - rocprofiler_context_id_t context_id, - rocprofiler_buffer_id_t buffer_id, - rocprofiler_dispatch_counting_service_callback_t callback, - void* callback_data_args) +rocprofiler_configure_buffer_dispatch_counting_service( + rocprofiler_context_id_t context_id, + rocprofiler_buffer_id_t buffer_id, + rocprofiler_dispatch_counting_service_cb_t callback, + void* callback_data_args) { return rocprofiler::counters::configure_buffered_dispatch( context_id, buffer_id, callback, callback_data_args); @@ -63,11 +63,11 @@ rocprofiler_configure_buffered_dispatch_counting_service( */ rocprofiler_status_t rocprofiler_configure_callback_dispatch_counting_service( - rocprofiler_context_id_t context_id, - 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_context_id_t context_id, + rocprofiler_dispatch_counting_service_cb_t dispatch_callback, + void* dispatch_callback_args, + rocprofiler_dispatch_counting_record_cb_t record_callback, + void* record_callback_args) { return rocprofiler::counters::configure_callback_dispatch(context_id, dispatch_callback, diff --git a/source/lib/rocprofiler-sdk/hip/abi.cpp b/source/lib/rocprofiler-sdk/hip/abi.cpp index 84ec610e26..112f175d03 100644 --- a/source/lib/rocprofiler-sdk/hip/abi.cpp +++ b/source/lib/rocprofiler-sdk/hip/abi.cpp @@ -20,8 +20,8 @@ // OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN // THE SOFTWARE. +#include #include -#include #include "lib/common/abi.hpp" #include "lib/common/defines.hpp" diff --git a/source/lib/rocprofiler-sdk/hip/details/format.hpp b/source/lib/rocprofiler-sdk/hip/details/format.hpp index 3bb58ac96c..8c8c74b05f 100644 --- a/source/lib/rocprofiler-sdk/hip/details/format.hpp +++ b/source/lib/rocprofiler-sdk/hip/details/format.hpp @@ -25,7 +25,6 @@ #include "lib/rocprofiler-sdk/hip/details/ostream.hpp" #include -#include #include // must be included after runtime api diff --git a/source/lib/rocprofiler-sdk/hip/stream.cpp b/source/lib/rocprofiler-sdk/hip/stream.cpp index c78d36d961..2b527f798c 100644 --- a/source/lib/rocprofiler-sdk/hip/stream.cpp +++ b/source/lib/rocprofiler-sdk/hip/stream.cpp @@ -169,7 +169,7 @@ FuncT create_write_functor(RetT (*func)(Args...)) return [](Args... args) -> RetT { using function_args_type = common::mpl::type_list; - using callback_api_data_t = rocprofiler_callback_tracing_stream_handle_data_t; + using callback_api_data_t = rocprofiler_callback_tracing_hip_stream_data_t; constexpr auto external_corr_id_domain_idx = hip_domain_info::external_correlation_id_domain_idx; @@ -179,8 +179,8 @@ FuncT create_write_functor(RetT (*func)(Args...)) auto buffered_contexts = tracing::buffered_context_data_vec_t{}; auto external_corr_ids = tracing::external_correlation_id_map_t{}; - tracing::populate_contexts(ROCPROFILER_CALLBACK_TRACING_HIP_STREAM_API, - ROCPROFILER_BUFFER_TRACING_HIP_STREAM_API, + tracing::populate_contexts(ROCPROFILER_CALLBACK_TRACING_HIP_STREAM, + ROCPROFILER_BUFFER_TRACING_HIP_STREAM, callback_contexts, buffered_contexts, external_corr_ids); @@ -208,7 +208,7 @@ FuncT create_write_functor(RetT (*func)(Args...)) internal_corr_id, external_corr_ids, ancestor_corr_id, - ROCPROFILER_CALLBACK_TRACING_HIP_STREAM_API, + ROCPROFILER_CALLBACK_TRACING_HIP_STREAM, ROCPROFILER_HIP_STREAM_CREATE, tracer_data); } @@ -232,7 +232,7 @@ FuncT create_destroy_functor(RetT (*func)(Args...)) using function_args_type = common::mpl::type_list; constexpr auto stream_idx = common::mpl::index_of::value; - using callback_api_data_t = rocprofiler_callback_tracing_stream_handle_data_t; + using callback_api_data_t = rocprofiler_callback_tracing_hip_stream_data_t; constexpr auto external_corr_id_domain_idx = hip_domain_info::external_correlation_id_domain_idx; @@ -242,8 +242,8 @@ FuncT create_destroy_functor(RetT (*func)(Args...)) auto buffered_contexts = tracing::buffered_context_data_vec_t{}; auto external_corr_ids = tracing::external_correlation_id_map_t{}; - tracing::populate_contexts(ROCPROFILER_CALLBACK_TRACING_HIP_STREAM_API, - ROCPROFILER_BUFFER_TRACING_HIP_STREAM_API, + tracing::populate_contexts(ROCPROFILER_CALLBACK_TRACING_HIP_STREAM, + ROCPROFILER_BUFFER_TRACING_HIP_STREAM, callback_contexts, buffered_contexts, external_corr_ids); @@ -268,7 +268,7 @@ FuncT create_destroy_functor(RetT (*func)(Args...)) internal_corr_id, external_corr_ids, ancestor_corr_id, - ROCPROFILER_CALLBACK_TRACING_HIP_STREAM_API, + ROCPROFILER_CALLBACK_TRACING_HIP_STREAM, ROCPROFILER_HIP_STREAM_DESTROY, tracer_data); } @@ -292,7 +292,7 @@ FuncT create_read_functor(RetT (*func)(Args...)) using function_args_type = common::mpl::type_list; constexpr auto stream_idx = common::mpl::index_of::value; - using callback_api_data_t = rocprofiler_callback_tracing_stream_handle_data_t; + using callback_api_data_t = rocprofiler_callback_tracing_hip_stream_data_t; constexpr auto external_corr_id_domain_idx = hip_domain_info::external_correlation_id_domain_idx; @@ -302,8 +302,8 @@ FuncT create_read_functor(RetT (*func)(Args...)) auto buffered_contexts = tracing::buffered_context_data_vec_t{}; auto external_corr_ids = tracing::external_correlation_id_map_t{}; - tracing::populate_contexts(ROCPROFILER_CALLBACK_TRACING_HIP_STREAM_API, - ROCPROFILER_BUFFER_TRACING_HIP_STREAM_API, + tracing::populate_contexts(ROCPROFILER_CALLBACK_TRACING_HIP_STREAM, + ROCPROFILER_BUFFER_TRACING_HIP_STREAM, callback_contexts, buffered_contexts, external_corr_ids); @@ -324,7 +324,7 @@ FuncT create_read_functor(RetT (*func)(Args...)) internal_corr_id, external_corr_ids, ancestor_corr_id, - ROCPROFILER_CALLBACK_TRACING_HIP_STREAM_API, + ROCPROFILER_CALLBACK_TRACING_HIP_STREAM, ROCPROFILER_HIP_STREAM_SET, tracer_data); } @@ -338,7 +338,7 @@ FuncT create_read_functor(RetT (*func)(Args...)) { tracing::execute_phase_exit_callbacks(callback_contexts, external_corr_ids, - ROCPROFILER_CALLBACK_TRACING_HIP_STREAM_API, + ROCPROFILER_CALLBACK_TRACING_HIP_STREAM, ROCPROFILER_HIP_STREAM_SET, tracer_data); } @@ -368,11 +368,11 @@ enable_stream_stack() if(itr->is_tracing_one_of(ROCPROFILER_CALLBACK_TRACING_MEMORY_COPY, ROCPROFILER_CALLBACK_TRACING_HIP_RUNTIME_API, ROCPROFILER_CALLBACK_TRACING_HIP_COMPILER_API, - ROCPROFILER_CALLBACK_TRACING_HIP_STREAM_API, + ROCPROFILER_CALLBACK_TRACING_HIP_STREAM, ROCPROFILER_BUFFER_TRACING_MEMORY_COPY, ROCPROFILER_BUFFER_TRACING_HIP_RUNTIME_API, ROCPROFILER_BUFFER_TRACING_HIP_COMPILER_API, - ROCPROFILER_BUFFER_TRACING_HIP_STREAM_API, + ROCPROFILER_BUFFER_TRACING_HIP_STREAM, ROCPROFILER_BUFFER_TRACING_HIP_RUNTIME_API_EXT, ROCPROFILER_BUFFER_TRACING_HIP_COMPILER_API_EXT)) return true; @@ -430,8 +430,8 @@ update_table(Tp* _orig, std::integral_constant) // index_of finds the first argument of that type. So find the first and last // arg of the given type and make sure it resolves to the same distance - assert(stream_idx == (num_args - rstream_idx - 1) && - "function has more than one stream argument"); + static_assert(stream_idx == (num_args - rstream_idx - 1), + "function has more than one stream argument"); // don't wrap the compiler API functions unless HIP compiler API tracing is enabled if constexpr(TableIdx == ROCPROFILER_HIP_TABLE_ID_Compiler) @@ -471,8 +471,8 @@ update_table(Tp* _orig, std::integral_constant) // index_of finds the first argument of that type. So find the first and last // arg of the given type and make sure it resolves to the same distance - assert(stream_idx == (num_args - rstream_idx - 1) && - "function has more than one stream argument"); + static_assert(stream_idx == (num_args - rstream_idx - 1), + "function has more than one stream argument"); // don't wrap the compiler API functions unless HIP compiler API tracing is enabled if constexpr(TableIdx == ROCPROFILER_HIP_TABLE_ID_Compiler) diff --git a/source/lib/rocprofiler-sdk/hip/utils.hpp b/source/lib/rocprofiler-sdk/hip/utils.hpp index 51be3f3dac..53916cbb5d 100644 --- a/source/lib/rocprofiler-sdk/hip/utils.hpp +++ b/source/lib/rocprofiler-sdk/hip/utils.hpp @@ -22,7 +22,7 @@ #pragma once -#include +#include #include "lib/common/mpl.hpp" #include "lib/common/stringize_arg.hpp" diff --git a/source/lib/rocprofiler-sdk/hsa/abi.cpp b/source/lib/rocprofiler-sdk/hsa/abi.cpp index dd12d14d81..b1dc079cdc 100644 --- a/source/lib/rocprofiler-sdk/hsa/abi.cpp +++ b/source/lib/rocprofiler-sdk/hsa/abi.cpp @@ -20,8 +20,8 @@ // OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN // THE SOFTWARE. +#include #include -#include #include "lib/common/abi.hpp" #include "lib/common/defines.hpp" diff --git a/source/lib/rocprofiler-sdk/hsa/queue.hpp b/source/lib/rocprofiler-sdk/hsa/queue.hpp index ed336f9263..e913d090e1 100644 --- a/source/lib/rocprofiler-sdk/hsa/queue.hpp +++ b/source/lib/rocprofiler-sdk/hsa/queue.hpp @@ -152,16 +152,16 @@ public: void set_state(queue_state state); private: - std::atomic _notifiers = {0}; - std::atomic _active_async_packets = {0}; - CoreApiTable _core_api = {}; - AmdExtTable _ext_api = {}; - const AgentCache& _agent; - rocprofiler::common::Synchronized _callbacks = {}; - hsa_queue_t* _intercept_queue = nullptr; - queue_state _state = queue_state::normal; - std::mutex _lock_queue; - hsa_signal_t _active_kernels = {.handle = 0}; + std::atomic _notifiers = {0}; + std::atomic _active_async_packets = {0}; + CoreApiTable _core_api = {}; + AmdExtTable _ext_api = {}; + const AgentCache& _agent; + common::Synchronized _callbacks = {}; + hsa_queue_t* _intercept_queue = nullptr; + queue_state _state = queue_state::normal; + std::mutex _lock_queue; + hsa_signal_t _active_kernels = {.handle = 0}; }; inline rocprofiler_queue_id_t diff --git a/source/lib/rocprofiler-sdk/hsa/scratch_memory.cpp b/source/lib/rocprofiler-sdk/hsa/scratch_memory.cpp index 441270aa53..75834fc241 100644 --- a/source/lib/rocprofiler-sdk/hsa/scratch_memory.cpp +++ b/source/lib/rocprofiler-sdk/hsa/scratch_memory.cpp @@ -444,6 +444,9 @@ impl(Args... args) constexpr auto external_corr_id_domain_idx = amd_tool_api_info::external_correlation_id_domain_idx; + // suppress unused-but-set-variable warning + common::consume_args(external_corr_id_domain_idx); + auto&& _tied_args = std::tie(args...); auto& event_data = std::get<0>(_tied_args); diff --git a/source/lib/rocprofiler-sdk/hsa/utils.hpp b/source/lib/rocprofiler-sdk/hsa/utils.hpp index 23f7c705c3..f4d76b89f7 100644 --- a/source/lib/rocprofiler-sdk/hsa/utils.hpp +++ b/source/lib/rocprofiler-sdk/hsa/utils.hpp @@ -22,7 +22,7 @@ #pragma once -#include +#include #include "lib/common/stringize_arg.hpp" diff --git a/source/lib/rocprofiler-sdk/marker/utils.hpp b/source/lib/rocprofiler-sdk/marker/utils.hpp index 50f7488519..c0dd168575 100644 --- a/source/lib/rocprofiler-sdk/marker/utils.hpp +++ b/source/lib/rocprofiler-sdk/marker/utils.hpp @@ -23,7 +23,7 @@ #pragma once #include -#include +#include #include "lib/common/mpl.hpp" #include "lib/common/stringize_arg.hpp" diff --git a/source/lib/rocprofiler-sdk/ompt/utils.hpp b/source/lib/rocprofiler-sdk/ompt/utils.hpp index b31bd8c00d..32da94c011 100644 --- a/source/lib/rocprofiler-sdk/ompt/utils.hpp +++ b/source/lib/rocprofiler-sdk/ompt/utils.hpp @@ -26,7 +26,6 @@ #include "lib/rocprofiler-sdk/ompt/details/format.hpp" #include -#include #include #include 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 a4bc4e7b9a..bdb284166a 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 @@ -207,7 +207,7 @@ record_callback(rocprofiler_dispatch_counting_service_data_t /*dispatch_data*/, void dispatch_callback(rocprofiler_dispatch_counting_service_data_t /*dispatch_data*/, - rocprofiler_profile_config_id_t* /*config*/, + rocprofiler_counter_config_id_t* /*config*/, rocprofiler_user_data_t* /*user_data*/, void* /*callback_data_args*/) {} @@ -215,7 +215,7 @@ dispatch_callback(rocprofiler_dispatch_counting_service_data_t /*dispatch_data*/ void set_profile(rocprofiler_context_id_t /*context_id*/, rocprofiler_agent_id_t /*agent*/, - rocprofiler_agent_set_profile_callback_t /*set_config*/, + rocprofiler_device_counting_agent_cb_t /*set_config*/, void*) {} diff --git a/source/lib/rocprofiler-sdk/rccl/abi.cpp b/source/lib/rocprofiler-sdk/rccl/abi.cpp index 91a53b6a77..bf516162b0 100644 --- a/source/lib/rocprofiler-sdk/rccl/abi.cpp +++ b/source/lib/rocprofiler-sdk/rccl/abi.cpp @@ -25,8 +25,8 @@ #include "lib/common/abi.hpp" #include "lib/common/defines.hpp" +#include #include -#include namespace rocprofiler { diff --git a/source/lib/rocprofiler-sdk/registration.cpp b/source/lib/rocprofiler-sdk/registration.cpp index 1027ea3fe4..d6a2ce7763 100644 --- a/source/lib/rocprofiler-sdk/registration.cpp +++ b/source/lib/rocprofiler-sdk/registration.cpp @@ -245,17 +245,19 @@ find_clients() return true; }; - auto emplace_client = [&data, priority_offset]( + constexpr auto client_id_size = sizeof(rocprofiler_client_id_t); + auto emplace_client = [&data, priority_offset]( std::string_view _name, void* _dlhandle, auto* _cfg_func) -> std::optional& { uint32_t _prio = priority_offset + data.size(); - return data.emplace_back(client_library{std::string{_name}, - _dlhandle, - _cfg_func, - nullptr, - rocprofiler_client_id_t{nullptr, _prio}, - rocprofiler_client_id_t{nullptr, _prio}}); + return data.emplace_back( + client_library{std::string{_name}, + _dlhandle, + _cfg_func, + nullptr, + rocprofiler_client_id_t{client_id_size, nullptr, _prio}, + rocprofiler_client_id_t{client_id_size, nullptr, _prio}}); }; auto rocprofiler_configure_dlsym = [](auto _handle) { diff --git a/source/lib/rocprofiler-sdk/rocdecode/abi.cpp b/source/lib/rocprofiler-sdk/rocdecode/abi.cpp index e60126f3d7..c02868eac2 100644 --- a/source/lib/rocprofiler-sdk/rocdecode/abi.cpp +++ b/source/lib/rocprofiler-sdk/rocdecode/abi.cpp @@ -25,8 +25,8 @@ #include "lib/common/abi.hpp" #include "lib/common/defines.hpp" +#include #include -#include namespace rocprofiler { diff --git a/source/lib/rocprofiler-sdk/rocjpeg/abi.cpp b/source/lib/rocprofiler-sdk/rocjpeg/abi.cpp index 4bdc016742..dd84d689ec 100644 --- a/source/lib/rocprofiler-sdk/rocjpeg/abi.cpp +++ b/source/lib/rocprofiler-sdk/rocjpeg/abi.cpp @@ -25,8 +25,8 @@ #include "lib/common/abi.hpp" #include "lib/common/defines.hpp" +#include #include -#include namespace rocprofiler { diff --git a/source/lib/rocprofiler-sdk/rocprofiler.cpp b/source/lib/rocprofiler-sdk/rocprofiler.cpp index 8fd30a2513..dfb27bf4f3 100644 --- a/source/lib/rocprofiler-sdk/rocprofiler.cpp +++ b/source/lib/rocprofiler-sdk/rocprofiler.cpp @@ -153,6 +153,15 @@ rocprofiler_get_version(uint32_t* major, uint32_t* minor, uint32_t* patch) return ROCPROFILER_STATUS_SUCCESS; } +rocprofiler_status_t +rocprofiler_get_version_triplet(rocprofiler_version_triplet_t* info) +{ + *info = {.major = ROCPROFILER_VERSION_MAJOR, + .minor = ROCPROFILER_VERSION_MINOR, + .patch = ROCPROFILER_VERSION_PATCH}; + return ROCPROFILER_STATUS_SUCCESS; +} + rocprofiler_status_t rocprofiler_get_timestamp(rocprofiler_timestamp_t* ts) { diff --git a/source/lib/rocprofiler-sdk/tests/hsa_barrier.cpp b/source/lib/rocprofiler-sdk/tests/hsa_barrier.cpp index 6a8afd5a70..76cc9dfd88 100644 --- a/source/lib/rocprofiler-sdk/tests/hsa_barrier.cpp +++ b/source/lib/rocprofiler-sdk/tests/hsa_barrier.cpp @@ -142,10 +142,10 @@ inject_barriers(hsa_barrier& barrier, QueueController::queue_map_t& queues) for(auto& [hsa_queue, fq] : queues) { - auto pkt = barrier.enqueue_packet(fq.get()); - ASSERT_EQ(pkt.has_value(), true); - hsa_barrier_and_packet_t* packets = (hsa_barrier_and_packet_t*) hsa_queue->base_address; - enqueue_pkt(hsa_queue, packets, pkt->barrier_and); + auto _pkt = barrier.enqueue_packet(fq.get()); + ASSERT_EQ(_pkt.has_value(), true); + hsa_barrier_and_packet_t* _packets = (hsa_barrier_and_packet_t*) hsa_queue->base_address; + enqueue_pkt(hsa_queue, _packets, _pkt->barrier_and); // Construct packet that will trigger async handler after barrier is released rocprofiler_packet post_barrier{}; @@ -158,7 +158,7 @@ inject_barriers(hsa_barrier& barrier, QueueController::queue_map_t& queues) 0, barrier_signal_handler, static_cast(completion_signal)); - enqueue_pkt(hsa_queue, packets, post_barrier.barrier_and); + enqueue_pkt(hsa_queue, _packets, post_barrier.barrier_and); } // Ensure that the barrier packet is reached on all queues diff --git a/source/lib/rocprofiler-sdk/tests/version.cpp b/source/lib/rocprofiler-sdk/tests/version.cpp index 79231e4aa8..65575229dc 100644 --- a/source/lib/rocprofiler-sdk/tests/version.cpp +++ b/source/lib/rocprofiler-sdk/tests/version.cpp @@ -20,6 +20,7 @@ // OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE // SOFTWARE. +#include #include #include #include @@ -28,6 +29,8 @@ #include +#include + TEST(rocprofiler_lib, version) { auto correct_version = std::tuple( diff --git a/source/libexec/rocprofiler-sdk/rocprofiler-avail/rocprofv3_avail.cpp b/source/libexec/rocprofiler-sdk/rocprofiler-avail/rocprofv3_avail.cpp index 1cafc5556e..008ac06a02 100644 --- a/source/libexec/rocprofiler-sdk/rocprofiler-avail/rocprofv3_avail.cpp +++ b/source/libexec/rocprofiler-sdk/rocprofiler-avail/rocprofv3_avail.cpp @@ -201,12 +201,15 @@ iterate_agent_counters_callback(rocprofiler_agent_id_t, auto* _counters_info = static_cast>*>(user_data); for(size_t i = 0; i < num_counters; i++) { - rocprofiler_counter_info_v0_t _info; + auto _info = rocprofiler_counter_info_v1_t{}; auto dimensions_data = std::vector{}; ROCPROFILER_CALL( - rocprofiler_iterate_counter_dimensions( - counters[i], dimensions_info_callback, static_cast(&dimensions_data)), - "iterate_dimension_info"); + rocprofiler_query_counter_info( + counters[i], ROCPROFILER_COUNTER_INFO_VERSION_1, static_cast(&_info)), + "Could not query counter_id"); + + dimensions_data = std::vector{ + _info.dimensions, _info.dimensions + _info.dimensions_count}; auto dimensions_info = std::vector>{}; dimensions_info.reserve(dimensions_data.size()); for(auto& dim : dimensions_data) @@ -220,10 +223,6 @@ iterate_agent_counters_callback(rocprofiler_agent_id_t, dimensions_info.emplace_back(dimensions); } counter_dim_info.emplace(counters[i].handle, dimensions_info); - ROCPROFILER_CALL( - rocprofiler_query_counter_info( - counters[i], ROCPROFILER_COUNTER_INFO_VERSION_0, static_cast(&_info)), - "Could not query counter_id"); auto counter = std::vector{}; diff --git a/tests/async-copy-tracing/validate.py b/tests/async-copy-tracing/validate.py index 14d7218f4c..6fd2f94e89 100644 --- a/tests/async-copy-tracing/validate.py +++ b/tests/async-copy-tracing/validate.py @@ -47,6 +47,14 @@ def get_operation(record, kind_name, op_name=None): return None +def get_operation_name(record, kind_idx, op_idx): + for idx, itr in enumerate(record["names"]): + if idx == kind_idx: + return itr["operations"][op_idx] + + return None + + def test_data_structure(input_data): """verify minimum amount of expected data is present""" data = input_data @@ -418,12 +426,6 @@ def test_ancestor_ids(input_data): len(memcopies) == num_hsa_memcopies ), "Expected number of memcopies to be same as number of async HSA (hsa_amd_memory_async_copy_on_engine) calls" - print(" > hip_memcopy_id : ", hip_memcopy_id) - print(" > has_async_memcopy_id : ", has_async_memcopy_id) - print( - " > len(hip_memcopies) : ", num_hsa_memcopies, len(hip_memcopies), len(memcopies) - ) - for tid in hip_memcopies: # We expect only 1 record with this internal id, per thread for corr_id, records in hip_memcopies[tid].items(): @@ -469,6 +471,8 @@ def test_ancestor_ids(input_data): def test_retired_correlation_ids(input_data): data = input_data sdk_data = data["rocprofiler-sdk-json-tool"] + buffer_records = sdk_data["buffer_records"] + api_name_info = {} def _sort_dict(inp): return dict(sorted(inp.items())) @@ -477,8 +481,13 @@ def test_retired_correlation_ids(input_data): for titr in ["hsa_api_traces", "marker_api_traces", "hip_api_traces"]: for itr in sdk_data["buffer_records"][titr]: corr_id = itr["correlation_id"]["internal"] + name = get_operation_name(buffer_records, itr["kind"], itr["operation"]) + assert corr_id not in api_corr_ids.keys() + assert name is not None, f"{itr}" + api_corr_ids[corr_id] = itr + api_name_info[corr_id] = name async_corr_ids = {} for titr in ["kernel_dispatch", "memory_copies"]: @@ -497,19 +506,55 @@ def test_retired_correlation_ids(input_data): async_corr_ids = _sort_dict(async_corr_ids) retired_corr_ids = _sort_dict(retired_corr_ids) - for cid, itr in async_corr_ids.items(): - assert cid in retired_corr_ids.keys() - retired_ts = retired_corr_ids[cid]["timestamp"] - end_ts = itr["end_timestamp"] - assert (retired_ts - end_ts) > 0, f"correlation-id: {cid}, data: {itr}" + # + # verify all the correlation ids were retired + # + num_api_corr_ids = len(api_corr_ids.keys()) + num_retired_corr_ids = len(retired_corr_ids.keys()) + missing_retired_corr_ids = [ + itr for itr in api_corr_ids.keys() if itr not in retired_corr_ids.keys() + ] + # log in case of failure + sys.stderr.flush() + for itr in missing_retired_corr_ids: + name = api_name_info[itr] + info = api_corr_ids[itr] + sys.stderr.write(f"- unretired corr id: {itr} :: {name} :: {info}\n") + sys.stderr.flush() + + assert ( + num_api_corr_ids == num_retired_corr_ids + ), f"correlation ids not retired:\n\t{missing_retired_corr_ids}" + + # + # verify the retirement timestamp is >= the end timestamp of the records + # for cid, itr in api_corr_ids.items(): assert cid in retired_corr_ids.keys() retired_ts = retired_corr_ids[cid]["timestamp"] end_ts = itr["end_timestamp"] - assert (retired_ts - end_ts) > 0, f"correlation-id: {cid}, data: {itr}" + name = api_name_info[cid] + assert ( + retired_ts - end_ts + ) >= 0, f"\n\tcorr: {cid}\n\tname: {name}\n\tdata: {itr}" - assert len(api_corr_ids.keys()) == (len(retired_corr_ids.keys())) + # allow the retired timestamp to be 10 usec earlier than async end timestamp + # since the async timestamps undergo conversion from the GPU clock domain to + # the CPU clock domain. 10 microseconds was arbitrarily chosen to be an + # acceptable amount of inaccuracy -- in an ideal world, retired_ts should + # always be >= end_ts + usec = 1000 + supported_fuzzing = 10 * usec + + for cid, itr in async_corr_ids.items(): + assert cid in retired_corr_ids.keys() + retired_ts = retired_corr_ids[cid]["timestamp"] + end_ts = itr["end_timestamp"] + name = api_name_info[cid] + assert ( + retired_ts - end_ts + ) >= -supported_fuzzing, f"\n\tcorr: {cid}\n\tname: {name}\n\tdata: {itr}" if __name__ == "__main__": diff --git a/tests/tools/json-tool.cpp b/tests/tools/json-tool.cpp index abec9da157..be37fefbb9 100644 --- a/tests/tools/json-tool.cpp +++ b/tests/tools/json-tool.cpp @@ -610,12 +610,12 @@ set_external_correlation_id(rocprofiler_thread_id_t t void dispatch_callback(rocprofiler_dispatch_counting_service_data_t dispatch_data, - rocprofiler_profile_config_id_t* config, + rocprofiler_counter_config_id_t* config, rocprofiler_user_data_t* /*user_data*/, void* /*callback_data_args*/) { static std::shared_mutex m_mutex = {}; - static std::unordered_map profile_cache = {}; + static std::unordered_map profile_cache = {}; auto search_cache = [&]() { if(auto pos = profile_cache.find(dispatch_data.dispatch_info.agent_id.handle); @@ -692,8 +692,8 @@ dispatch_callback(rocprofiler_dispatch_counting_service_data_t dispatch_data, } // Create a colleciton profile for the counters - rocprofiler_profile_config_id_t profile = {.handle = 0}; - ROCPROFILER_CALL(rocprofiler_create_profile_config(dispatch_data.dispatch_info.agent_id, + rocprofiler_counter_config_id_t profile = {.handle = 0}; + ROCPROFILER_CALL(rocprofiler_create_counter_config(dispatch_data.dispatch_info.agent_id, collect_counters.data(), collect_counters.size(), &profile), @@ -1682,7 +1682,7 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data) "buffer tracing service for ompt configure"); ROCPROFILER_CALL( - rocprofiler_configure_buffered_dispatch_counting_service( + rocprofiler_configure_buffer_dispatch_counting_service( counter_collection_ctx, counter_collection_buffer, dispatch_callback, nullptr), "setup buffered service");