From 957bb7a4e55176d92887e7296f2d3d8ba1c7f325 Mon Sep 17 00:00:00 2001 From: "Jonathan R. Madsen" Date: Wed, 22 May 2024 15:51:12 -0500 Subject: [PATCH] Adding Perfetto support (#867) * Perfetto submodule * include/rocprofiler-sdk/cxx/perfetto.hpp - adapted from tests/common/perfetto.hpp - updated json-tool to use * Update include/rocprofiler-sdk/cxx - add details/delimit.hpp - add details/join.hpp - extend details/mpl.hpp - extend details/operators.hpp * Update lib/rocprofiler-sdk/hsa/async_copy.cpp - update MEMORY_COPY direction names * Preliminary perfetto support * Update lib/rocprofiler-sdk-tool/generatePerfetto.cpp - fix getting roctx msg vs. buffer operation name * Temporary variable restructuring * Perfetto patches after rebasing onto main * Revert lib/rocprofiler-sdk/hsa/async_copy.cpp - revert name * Update lib/rocprofiler-sdk-tool/generatePerfetto.cpp - fix ReadTrace * Update tests/bin/hip-in-libraries - sleep_for * Support PFTRACE output format option in rocprofv3 * Change perfetto logging * Update rocprofv3 tests to generate pftrace output * Minor tweak to json-tool.cpp * Update requirements.txt for perfetto testing * Fix data race on amount_read in generatePerfetto.cpp * Add testing for pftrace output - relatively simple testing which verifies that the pftrace file has the same number of entries as JSON data for HIP/HSA/marker/kernel/memory_copy * Fix import in perfetto_reader.py * Fix data race in generatePerfetto.cpp --- .gitmodules | 3 + external/CMakeLists.txt | 22 + external/perfetto | 1 + requirements.txt | 3 + source/bin/rocprofv3.sh | 7 +- .../rocprofiler-sdk/cxx/CMakeLists.txt | 3 +- .../cxx/details/CMakeLists.txt | 2 +- .../rocprofiler-sdk/cxx/details/delimit.hpp | 158 ++++ .../rocprofiler-sdk/cxx/details/join.hpp | 281 +++++++ .../rocprofiler-sdk/cxx/details/mpl.hpp | 173 ++++ .../include/rocprofiler-sdk/cxx/operators.hpp | 74 +- .../include/rocprofiler-sdk/cxx/perfetto.hpp | 168 ++++ .../lib/rocprofiler-sdk-tool/CMakeLists.txt | 12 +- source/lib/rocprofiler-sdk-tool/config.cpp | 7 +- source/lib/rocprofiler-sdk-tool/config.hpp | 1 + .../lib/rocprofiler-sdk-tool/generateCSV.cpp | 1 + .../rocprofiler-sdk-tool/generatePerfetto.cpp | 529 ++++++++++++ .../rocprofiler-sdk-tool/generatePerfetto.hpp | 45 ++ source/lib/rocprofiler-sdk-tool/tool.cpp | 14 + .../bin/hip-in-libraries/hip-in-libraries.cpp | 5 +- tests/common/CMakeLists.txt | 52 +- tests/common/perfetto.hpp | 202 +---- tests/pytest-packages/CMakeLists.txt | 9 +- .../pytest_utils/perfetto_reader.py | 750 ++++++++++++++++++ tests/pytest-packages/tests/__init__.py | 23 + tests/pytest-packages/tests/rocprofv3.py | 48 ++ .../hsa-queue-dependency/CMakeLists.txt | 6 +- .../hsa-queue-dependency/conftest.py | 12 + .../hsa-queue-dependency/validate.py | 8 + .../tracing-hip-in-libraries/CMakeLists.txt | 33 +- .../tracing-hip-in-libraries/conftest.py | 12 + .../tracing-hip-in-libraries/validate.py | 9 + .../rocprofv3/tracing-plus-cc/CMakeLists.txt | 4 +- tests/rocprofv3/tracing-plus-cc/conftest.py | 12 + tests/rocprofv3/tracing-plus-cc/validate.py | 8 + tests/rocprofv3/tracing/CMakeLists.txt | 24 +- tests/rocprofv3/tracing/conftest.py | 12 + tests/rocprofv3/tracing/validate.py | 8 + tests/tools/json-tool.cpp | 80 +- 39 files changed, 2530 insertions(+), 291 deletions(-) create mode 160000 external/perfetto create mode 100644 source/include/rocprofiler-sdk/cxx/details/delimit.hpp create mode 100644 source/include/rocprofiler-sdk/cxx/details/join.hpp create mode 100644 source/include/rocprofiler-sdk/cxx/perfetto.hpp create mode 100644 source/lib/rocprofiler-sdk-tool/generatePerfetto.cpp create mode 100644 source/lib/rocprofiler-sdk-tool/generatePerfetto.hpp create mode 100644 tests/pytest-packages/pytest_utils/perfetto_reader.py create mode 100644 tests/pytest-packages/tests/__init__.py create mode 100644 tests/pytest-packages/tests/rocprofv3.py diff --git a/.gitmodules b/.gitmodules index 1ad4746ab1..4985459020 100644 --- a/.gitmodules +++ b/.gitmodules @@ -19,3 +19,6 @@ [submodule "external/filesystem"] path = external/filesystem url = https://github.com/gulrak/filesystem.git +[submodule "external/perfetto"] + path = external/perfetto + url = https://android.googlesource.com/platform/external/perfetto diff --git a/external/CMakeLists.txt b/external/CMakeLists.txt index c62a4d3a0e..0289cce2ee 100644 --- a/external/CMakeLists.txt +++ b/external/CMakeLists.txt @@ -150,3 +150,25 @@ if(ROCPROFILER_BUILD_DOCS) REPO_URL https://github.com/jothepro/doxygen-awesome-css.git REPO_BRANCH "main") endif() + +rocprofiler_checkout_git_submodule( + RECURSIVE + RELATIVE_PATH external/perfetto + WORKING_DIRECTORY ${PROJECT_SOURCE_DIR} + TEST_FILE meson.build + REPO_URL https://android.googlesource.com/platform/external/perfetto external/perfetto + REPO_BRANCH "v44.0") + +add_library(rocprofiler-perfetto-static-library STATIC) +target_sources( + rocprofiler-perfetto-static-library + PRIVATE ${PROJECT_SOURCE_DIR}/external/perfetto/sdk/perfetto.h + ${PROJECT_SOURCE_DIR}/external/perfetto/sdk/perfetto.cc) +target_include_directories( + rocprofiler-perfetto-static-library SYSTEM + INTERFACE $) +set_target_properties( + rocprofiler-perfetto-static-library PROPERTIES POSITION_INDEPENDENT_CODE ON + OUTPUT_NAME rocprofiler-sdk-perfetto) +target_link_libraries(rocprofiler-perfetto + INTERFACE $) diff --git a/external/perfetto b/external/perfetto new file mode 160000 index 0000000000..eb5ef24c58 --- /dev/null +++ b/external/perfetto @@ -0,0 +1 @@ +Subproject commit eb5ef24c58d13cec289d733d03f0f3f0ed321b12 diff --git a/requirements.txt b/requirements.txt index 169854cc50..4195d9d00f 100644 --- a/requirements.txt +++ b/requirements.txt @@ -3,7 +3,10 @@ clang-format>=11.0.0,<12.0.0 clang-tidy>=15.0.0,<18.0.0 cmake>=3.21.0 cmake-format +dataclasses +numpy pandas +perfetto pycobertura pytest pyyaml diff --git a/source/bin/rocprofv3.sh b/source/bin/rocprofv3.sh index 692dcda182..2707f3321a 100755 --- a/source/bin/rocprofv3.sh +++ b/source/bin/rocprofv3.sh @@ -55,9 +55,14 @@ usage() { echo -e "\t#${GREY} usage (with custom dir): rocprofv3 --hsa-trace -d -o ${RESET}\n" echo -e "" echo -e "${GREEN}-d | --output-directory ${RESET} For adding output path where the output files will be saved" - echo -e "${GREEN} | --output-format ${RESET} For adding output format" echo -e "\t#${GREY} usage (with custom dir): rocprofv3 --hsa-trace -d ${RESET}" echo -e "" + echo -e "${GREEN}--output-format ${RESET} For specifying output format. Case-insensitive, comma separated. Options: CSV, JSON, PFTRACE" + echo -e "\t#${GREY} Examples:" + echo -e "\t#${GREY} (JSON output): rocprofv3 --sys-trace --output-format JSON " + echo -e "\t#${GREY} (JSON + CSV): rocprofv3 --sys-trace --output-format JSON,CSV " + echo -e "\t#${GREY} (JSON + PFTRACE): rocprofv3 --sys-trace --output-format JSON,PFTRACE " + echo -e "" echo -e "${GREEN}-M | --mangled-kernels ${RESET} Do not demangle the kernel names" echo -e "${GREEN}-T | --truncate-kernels ${RESET} Truncate the demangled kernel names" echo -e "" diff --git a/source/include/rocprofiler-sdk/cxx/CMakeLists.txt b/source/include/rocprofiler-sdk/cxx/CMakeLists.txt index 9e3b92c117..da9ddc76e1 100644 --- a/source/include/rocprofiler-sdk/cxx/CMakeLists.txt +++ b/source/include/rocprofiler-sdk/cxx/CMakeLists.txt @@ -3,7 +3,8 @@ # Installation of public C++ headers # # -set(ROCPROFILER_CXX_HEADER_FILES hash.hpp name_info.hpp operators.hpp serialization.hpp) +set(ROCPROFILER_CXX_HEADER_FILES hash.hpp name_info.hpp operators.hpp perfetto.hpp + serialization.hpp) install( FILES ${ROCPROFILER_CXX_HEADER_FILES} diff --git a/source/include/rocprofiler-sdk/cxx/details/CMakeLists.txt b/source/include/rocprofiler-sdk/cxx/details/CMakeLists.txt index bdb3729883..d8025a1305 100644 --- a/source/include/rocprofiler-sdk/cxx/details/CMakeLists.txt +++ b/source/include/rocprofiler-sdk/cxx/details/CMakeLists.txt @@ -3,7 +3,7 @@ # Installation of public C++ headers (implementations) # # -set(ROCPROFILER_CXX_DETAILS_HEADER_FILES mpl.hpp name_info.hpp) +set(ROCPROFILER_CXX_DETAILS_HEADER_FILES delimit.hpp join.hpp mpl.hpp name_info.hpp) install( FILES ${ROCPROFILER_CXX_DETAILS_HEADER_FILES} diff --git a/source/include/rocprofiler-sdk/cxx/details/delimit.hpp b/source/include/rocprofiler-sdk/cxx/details/delimit.hpp new file mode 100644 index 0000000000..d660b515f9 --- /dev/null +++ b/source/include/rocprofiler-sdk/cxx/details/delimit.hpp @@ -0,0 +1,158 @@ +// MIT License +// +// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + +#pragma once + +#include + +#include +#include +#include +#include +#include + +namespace rocprofiler +{ +namespace sdk +{ +namespace parse +{ +template +inline Tp +from_string(const std::string& str) +{ + auto ss = std::stringstream{str}; + auto val = Tp{}; + ss >> val; + return val; +} + +template +inline Tp +from_string(const char* cstr) +{ + auto ss = std::stringstream{cstr}; + auto val = Tp{}; + ss >> val; + return val; +} + +/// \brief tokenize a string into a set +/// +template , + typename ValueT = typename ContainerT::value_type, + typename PredicateT = std::function> +inline ContainerT +tokenize( + std::string_view line, + std::string_view delimiters = "\"',;: ", + PredicateT&& predicate = [](ValueT&& s) -> ValueT { return s; }) +{ + using value_type = ValueT; + + size_t _beginp = 0; // position that is the beginning of the new string + size_t _delimp = 0; // position of the delimiter in the string + ContainerT _result = {}; + if(mpl::reserve(_result, 0)) + { + size_t _nmax = 0; + for(char itr : line) + { + if(delimiters.find(itr) != std::string::npos) ++_nmax; + } + mpl::reserve(_result, _nmax); + } + while(_beginp < line.length() && _delimp < line.length()) + { + // find the first character (starting at _delimp) that is not a delimiter + _beginp = line.find_first_not_of(delimiters, _delimp); + // if no a character after or at _end that is not a delimiter is not found + // then we are done + if(_beginp == std::string::npos) break; + // starting at the position of the new string, find the next delimiter + _delimp = line.find_first_of(delimiters, _beginp); + + auto _tmp = value_type{}; + // starting at the position of the new string, get the characters + // between this position and the next delimiter + if(_beginp < line.length()) _tmp = line.substr(_beginp, _delimp - _beginp); + + // don't add empty strings + if(!_tmp.empty()) + { + mpl::emplace(_result, std::forward(predicate)(std::move(_tmp))); + } + } + return _result; +} + +/// \brief apply a string transformation to substring in between a common delimiter. +/// +template > +inline std::string +str_transform(std::string_view input, + std::string_view _begin, + std::string_view _end, + PredicateT&& predicate) +{ + size_t _beg_pos = 0; // position that is the beginning of the new string + size_t _end_pos = 0; // position of the delimiter in the string + std::string _result = std::string{input}; + while(_beg_pos < _result.length() && _end_pos < _result.length()) + { + // find the first sequence of characters after the end-position + _beg_pos = _result.find(_begin, _end_pos); + + // if sequence wasn't found, we are done + if(_beg_pos == std::string::npos) break; + + // starting after the position of the first delimiter, find the end sequence + if(!_end.empty()) + _end_pos = _result.find(_end, _beg_pos + 1); + else + _end_pos = _beg_pos + _begin.length(); + + // break if not found + if(_end_pos == std::string::npos) break; + + // length of the substr being operated on + auto _len = _end_pos - _beg_pos; + + // get the substring between the two delimiters (including first delimiter) + auto _sub = _result.substr(_beg_pos, _len); + + // apply the transform + auto _transformed = predicate(_sub); + + // only replace if necessary + if(_sub != _transformed) + { + _result = _result.replace(_beg_pos, _len, _transformed); + // move end to the end of transformed string + _end_pos = _beg_pos + _transformed.length(); + } + } + return _result; +} +} // namespace parse +} // namespace sdk +} // namespace rocprofiler diff --git a/source/include/rocprofiler-sdk/cxx/details/join.hpp b/source/include/rocprofiler-sdk/cxx/details/join.hpp new file mode 100644 index 0000000000..58310a6e62 --- /dev/null +++ b/source/include/rocprofiler-sdk/cxx/details/join.hpp @@ -0,0 +1,281 @@ +// MIT License +// +// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + +#pragma once + +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace rocprofiler +{ +namespace sdk +{ +namespace join +{ +template +inline void +consume_args(ArgsT&&...) +{} + +enum +{ + NoQuoteStrings = 0x0, + QuoteStrings = 0x1 +}; + +template +struct triplet_config +{ + static constexpr auto index() { return Idx; } + std::string_view delimiter = {}; + std::string_view prefix = {}; + std::string_view suffix = {}; +}; + +using generic_config = triplet_config<0>; +using array_config = triplet_config<1>; +using pair_config = triplet_config<2>; + +struct config : generic_config +{ + using format_flags_t = std::ios_base::fmtflags; + using base_type = generic_config; + + config() = default; + ~config() = default; + config(const config&) = default; + config(config&&) noexcept = default; + + config& operator=(const config&) = default; + config& operator=(config&&) noexcept = default; + + // converting constructor + config(std::string_view _delim) + : base_type{_delim} + {} + + // converting constructor + config(const char* const _delim) + : base_type{_delim} + {} + + config(generic_config _cfg) + : base_type{_cfg} + {} + + config(array_config _cfg) + : array{_cfg} + {} + + config(pair_config _cfg) + : pair{_cfg} + {} + + config(generic_config _generic, array_config _array) + : base_type{_generic} + , array{_array} + {} + + config(generic_config _generic, pair_config _pair) + : base_type{_generic} + , pair{_pair} + {} + + config(array_config _array, pair_config _pair) + : array{_array} + , pair{_pair} + {} + + format_flags_t flags = std::ios_base::boolalpha; + array_config array = {", ", "[", "]"}; + pair_config pair = {", ", "{", "}"}; +}; + +namespace impl +{ +template +inline decltype(auto) +join_arg(config _cfg, ArgT&& _v) +{ + using arg_type = mpl::basic_identity_t; + + constexpr bool _is_string_type = mpl::is_string_type::value; + constexpr bool _is_iterable = mpl::is_iterable(0); + constexpr bool _has_traits_type = mpl::has_traits(0); + constexpr bool _has_key_type = mpl::has_key_type(0); + constexpr bool _has_value_type = mpl::has_value_type(0); + constexpr bool _has_mapped_type = mpl::has_mapped_type(0); + + if constexpr(_is_string_type) + { + if constexpr(TraitT == QuoteStrings) + { + return std::string{"\""} + std::string{std::forward(_v)} + std::string{"\""}; + } + else + { + return std::forward(_v); + } + } + else if constexpr(_is_iterable && !_has_traits_type && + (_has_value_type || (_has_key_type && _has_mapped_type))) + { + if constexpr(_has_key_type && _has_mapped_type) + { + std::stringstream _ss{}; + _ss.setf(_cfg.flags); + for(auto&& itr : std::forward(_v)) + _ss << _cfg.array.delimiter << _cfg.pair.prefix << join_arg(_cfg, itr.first) + << _cfg.pair.delimiter << join_arg(_cfg, itr.second) + << _cfg.pair.suffix; + auto _ret = _ss.str(); + auto&& _len = _cfg.array.delimiter.length(); + return (_ret.length() > _len) ? (std::string{_cfg.array.prefix} + _ret.substr(_len) + + std::string{_cfg.array.suffix}) + : std::string{}; + } + else if constexpr(_has_value_type) + { + std::stringstream _ss{}; + _ss.setf(_cfg.flags); + for(auto&& itr : std::forward(_v)) + _ss << _cfg.array.delimiter << join_arg(_cfg, itr); + auto _ret = _ss.str(); + auto&& _len = _cfg.array.delimiter.length(); + return (_ret.length() > _len) ? (std::string{_cfg.array.prefix} + _ret.substr(_len) + + std::string{_cfg.array.suffix}) + : std::string{}; + } + } + else if constexpr(mpl::supports_ostream(0)) + { + return std::forward(_v); + } + else + { + static_assert(_is_iterable, "Type is not iterable"); + static_assert(!_has_traits_type, "Type has a traits type"); + if constexpr(!_has_value_type) + { + static_assert(_has_key_type && _has_mapped_type, + "Type must have a key_type and mapped_type if there is no value_type"); + } + else + { + static_assert(_has_value_type, + "Type must have a value_type if there is no key_type and mapped_type"); + } + static_assert(std::is_empty::value, + "Error! argument type cannot be written to output stream"); + } + // suppress any unused but set variable warnings + consume_args(_is_string_type, + _is_iterable, + _has_traits_type, + _has_key_type, + _has_value_type, + _has_mapped_type); +} +} // namespace impl + +template +auto +join(config _cfg, Args&&... _args) +{ + static_assert(std::is_trivially_copyable::value, + "Error! config is not trivially copyable"); + + std::stringstream _ss{}; + _ss.setf(_cfg.flags); + ((_ss << ((mpl::is_empty(_args)) ? std::string_view{} : std::string_view{_cfg.delimiter}) + << impl::join_arg(_cfg, _args)), + ...); + auto _ret = _ss.str(); + auto&& _len = _cfg.delimiter.length(); + auto _cmp = + strncmp(std::string_view{_ret}.data(), std::string_view{_cfg.delimiter}.data(), _len) == 0; + return (_ret.length() > _len) ? (std::string{_cfg.prefix} + + ((_cmp) ? _ret.substr(_len) : _ret) + std::string{_cfg.suffix}) + : std::string{}; +} + +template +auto +join(std::array&& _delims, Args&&... _args) +{ + auto _cfg = config{}; + _cfg.delimiter = _delims.at(0); + _cfg.prefix = _delims.at(1); + _cfg.suffix = _delims.at(2); + return join(_cfg, std::forward(_args)...); +} + +template ::value, int> = 0> +auto +join(DelimT&& _delim, Args&&... _args) +{ + using delim_type = mpl::basic_identity_t; + + if constexpr(std::is_constructible::value) + { + auto _cfg = config{std::forward(_delim)}; + return join(_cfg, std::forward(_args)...); + } + else if constexpr(std::is_same::value) + { + auto _cfg = config{}; + const char _delim_c[2] = {_delim, '\0'}; + _cfg.delimiter = _delim_c; + return join(_cfg, std::forward(_args)...); + } + else + { + auto _cfg = config{}; + _cfg.delimiter = std::string_view{_delim}; + return join(_cfg, std::forward(_args)...); + } +} + +template +auto +quoted(ArgT&& _arg) +{ + auto _cfg = config{}; + _cfg.prefix = "\""; + _cfg.suffix = "\""; + return join(_cfg, std::forward(_arg)); +} +} // namespace join +} // namespace sdk +} // namespace rocprofiler diff --git a/source/include/rocprofiler-sdk/cxx/details/mpl.hpp b/source/include/rocprofiler-sdk/cxx/details/mpl.hpp index 56bd4ee884..0299fe72a1 100644 --- a/source/include/rocprofiler-sdk/cxx/details/mpl.hpp +++ b/source/include/rocprofiler-sdk/cxx/details/mpl.hpp @@ -23,17 +23,61 @@ #pragma once +#include +#include +#include #include #include #include #include +#define ROCPROFILER_IMPL_HAS_CONCEPT(NAME, TRAIT) \ + template \ + inline constexpr bool NAME(int) \ + { \ + return true; \ + } \ + \ + template \ + inline constexpr bool NAME(long) \ + { \ + return false; \ + } + +#define ROCPROFILER_IMPL_SFINAE_CONCEPT(NAME, ...) \ + template \ + struct NAME \ + { \ + private: \ + static constexpr auto sfinae(int) -> decltype(__VA_ARGS__, bool()) { return true; } \ + \ + static constexpr auto sfinae(long) { return false; } \ + \ + public: \ + static constexpr bool value = sfinae(0); \ + constexpr auto operator()() const { return sfinae(0); } \ + }; + namespace rocprofiler { namespace sdk { namespace mpl { +template +struct unqualified_identity +{ + using type = std::remove_cv_t>>; +}; + +template +using unqualified_identity_t = typename unqualified_identity::type; + +template +struct is_same_unqualified_identity +: std::is_same, unqualified_identity_t> +{}; + template struct string_support { @@ -79,6 +123,135 @@ struct string_support type operator()(const char* val) const { return type{val}; } }; + +namespace impl +{ +template +struct is_string_type : std::false_type +{}; + +template <> +struct is_string_type : std::true_type +{}; + +template <> +struct is_string_type : std::true_type +{}; + +template <> +struct is_string_type : std::true_type +{}; + +template <> +struct is_string_type : std::true_type +{}; +} // namespace impl + +template +struct is_string_type : impl::is_string_type> +{}; + +// template +// struct can_stringify +// { +// private: +// static constexpr auto sfinae(int) +// -> decltype(std::declval() << std::declval(), bool()) +// { +// return true; +// } + +// static constexpr auto sfinae(long) { return false; } + +// public: +// static constexpr bool value = sfinae(0); +// constexpr auto operator()() const { return sfinae(0); } +// }; + +ROCPROFILER_IMPL_HAS_CONCEPT(has_traits, traits_type) +ROCPROFILER_IMPL_HAS_CONCEPT(has_value_type, value_type) +ROCPROFILER_IMPL_HAS_CONCEPT(has_key_type, key_type) +ROCPROFILER_IMPL_HAS_CONCEPT(has_mapped_type, mapped_type) + +ROCPROFILER_IMPL_SFINAE_CONCEPT(has_empty_member_function, std::declval().empty()) +ROCPROFILER_IMPL_SFINAE_CONCEPT(can_stringify, std::declval() << std::declval()) +ROCPROFILER_IMPL_SFINAE_CONCEPT(is_iterable, + std::begin(std::declval()), + std::end(std::declval())) + +// compatability +template +using supports_ostream = can_stringify; + +template +inline bool +is_empty(ArgT&& _v) +{ + using arg_type = unqualified_identity_t; + + if constexpr(has_empty_member_function::value) + { + return std::forward(_v).empty(); + } + else if constexpr(is_string_type::value) + { + static_assert(std::is_constructible::value, + "not string_view constructible"); + return std::string_view{std::forward(_v)}.empty(); + } + + return false; +} + +namespace impl +{ +template +inline auto +emplace(ContainerT& _c, int, Args&&... _args) + -> decltype(_c.emplace_back(std::forward(_args)...)) +{ + return _c.emplace_back(std::forward(_args)...); +} + +template +inline auto +emplace(ContainerT& _c, long, Args&&... _args) -> decltype(_c.emplace(std::forward(_args)...)) +{ + return _c.emplace(std::forward(_args)...); +} + +template +inline auto +reserve(ContainerT& _c, int, ArgT _arg) -> decltype(_c.reserve(_arg), bool()) +{ + _c.reserve(_arg); + return true; +} + +template +inline auto +reserve(ContainerT&, long, ArgT) +{ + return false; +} +} // namespace impl + +template +inline auto +emplace(ContainerT& _c, Args&&... _args) +{ + return impl::emplace(_c, 0, std::forward(_args)...); +} + +template +inline auto +reserve(ContainerT& _c, ArgT _arg) +{ + return impl::reserve(_c, 0, _arg); +} } // namespace mpl } // namespace sdk } // namespace rocprofiler + +#undef ROCPROFILER_IMPL_HAS_CONCEPT +#undef ROCPROFILER_IMPL_SFINAE_CONCEPT diff --git a/source/include/rocprofiler-sdk/cxx/operators.hpp b/source/include/rocprofiler-sdk/cxx/operators.hpp index dd4bcc6dac..dd8f59a71b 100644 --- a/source/include/rocprofiler-sdk/cxx/operators.hpp +++ b/source/include/rocprofiler-sdk/cxx/operators.hpp @@ -32,7 +32,11 @@ #define ROCPROFILER_CXX_DECLARE_OPERATORS(TYPE) \ bool operator==(TYPE lhs, TYPE rhs) ROCPROFILER_ATTRIBUTE(pure); \ - bool operator!=(TYPE lhs, TYPE rhs) ROCPROFILER_ATTRIBUTE(pure); + bool operator!=(TYPE lhs, TYPE rhs) ROCPROFILER_ATTRIBUTE(pure); \ + bool operator<(TYPE lhs, TYPE rhs) ROCPROFILER_ATTRIBUTE(pure); \ + bool operator>(TYPE lhs, TYPE rhs) ROCPROFILER_ATTRIBUTE(pure); \ + bool operator<=(TYPE lhs, TYPE rhs) ROCPROFILER_ATTRIBUTE(pure); \ + bool operator>=(TYPE lhs, TYPE rhs) ROCPROFILER_ATTRIBUTE(pure); #define ROCPROFILER_CXX_DEFINE_NE_OPERATOR(TYPE) \ inline bool operator!=(TYPE lhs, TYPE rhs) { return !(lhs == rhs); } @@ -43,6 +47,17 @@ return ::rocprofiler::sdk::operators::equal(lhs, rhs); \ } +#define ROCPROFILER_CXX_DEFINE_LT_HANDLE_OPERATOR(TYPE) \ + inline bool operator<(TYPE lhs, TYPE rhs) \ + { \ + return ::rocprofiler::sdk::operators::less(lhs, rhs); \ + } + +#define ROCPROFILER_CXX_DEFINE_COMPARE_OPERATORS(TYPE) \ + inline bool operator>(TYPE lhs, TYPE rhs) { return (lhs == rhs || !(lhs < rhs)); } \ + inline bool operator<=(TYPE lhs, TYPE rhs) { return (lhs == rhs || lhs < rhs); } \ + inline bool operator>=(TYPE lhs, TYPE rhs) { return !(lhs < rhs); } + namespace rocprofiler { namespace sdk @@ -53,6 +68,10 @@ template bool equal(Tp lhs, Tp rhs) ROCPROFILER_ATTRIBUTE(pure); +template +bool +less(Tp lhs, Tp rhs) ROCPROFILER_ATTRIBUTE(pure); + template bool equal(Tp lhs, Tp rhs) @@ -60,6 +79,14 @@ equal(Tp lhs, Tp rhs) static_assert(sizeof(Tp) == sizeof(uint64_t), "error! only for opaque handle types"); return lhs.handle == rhs.handle; } + +template +bool +less(Tp lhs, Tp rhs) +{ + static_assert(sizeof(Tp) == sizeof(uint64_t), "error! only for opaque handle types"); + return lhs.handle < rhs.handle; +} } // namespace operators } // namespace sdk } // namespace rocprofiler @@ -116,7 +143,52 @@ ROCPROFILER_CXX_DEFINE_NE_OPERATOR(hsa_executable_t) ROCPROFILER_CXX_DEFINE_NE_OPERATOR(const rocprofiler_agent_v0_t&) ROCPROFILER_CXX_DEFINE_NE_OPERATOR(rocprofiler_dim3_t) +// definitions of operator< +ROCPROFILER_CXX_DEFINE_LT_HANDLE_OPERATOR(rocprofiler_context_id_t) +ROCPROFILER_CXX_DEFINE_LT_HANDLE_OPERATOR(rocprofiler_agent_id_t) +ROCPROFILER_CXX_DEFINE_LT_HANDLE_OPERATOR(rocprofiler_queue_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_callback_thread_t) +ROCPROFILER_CXX_DEFINE_LT_HANDLE_OPERATOR(hsa_agent_t) +ROCPROFILER_CXX_DEFINE_LT_HANDLE_OPERATOR(hsa_signal_t) +ROCPROFILER_CXX_DEFINE_LT_HANDLE_OPERATOR(hsa_executable_t) + +inline bool +operator<(const rocprofiler_agent_v0_t& lhs, const rocprofiler_agent_v0_t& rhs) +{ + return (lhs.id < rhs.id); +} + +inline bool +operator<(rocprofiler_dim3_t lhs, rocprofiler_dim3_t rhs) +{ + const auto magnitude = [](rocprofiler_dim3_t dim_v) { return dim_v.x * dim_v.y * dim_v.z; }; + auto lhs_m = magnitude(lhs); + auto rhs_m = magnitude(rhs); + + return (lhs_m == rhs_m) ? std::tie(lhs.x, lhs.y, lhs.z) < std::tie(rhs.x, rhs.y, rhs.z) + : (lhs_m < rhs_m); +} + +// definitions of operator>, operator<=, operator>= +ROCPROFILER_CXX_DEFINE_COMPARE_OPERATORS(rocprofiler_context_id_t) +ROCPROFILER_CXX_DEFINE_COMPARE_OPERATORS(rocprofiler_agent_id_t) +ROCPROFILER_CXX_DEFINE_COMPARE_OPERATORS(rocprofiler_queue_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_callback_thread_t) +ROCPROFILER_CXX_DEFINE_COMPARE_OPERATORS(hsa_agent_t) +ROCPROFILER_CXX_DEFINE_COMPARE_OPERATORS(hsa_signal_t) +ROCPROFILER_CXX_DEFINE_COMPARE_OPERATORS(hsa_executable_t) +ROCPROFILER_CXX_DEFINE_COMPARE_OPERATORS(const rocprofiler_agent_v0_t&) +ROCPROFILER_CXX_DEFINE_COMPARE_OPERATORS(rocprofiler_dim3_t) + // cleanup defines #undef ROCPROFILER_CXX_DECLARE_OPERATORS #undef ROCPROFILER_CXX_DEFINE_NE_OPERATOR #undef ROCPROFILER_CXX_DEFINE_EQ_HANDLE_OPERATOR +#undef ROCPROFILER_CXX_DEFINE_LT_HANDLE_OPERATOR +#undef ROCPROFILER_CXX_DEFINE_COMPARE_OPERATORS diff --git a/source/include/rocprofiler-sdk/cxx/perfetto.hpp b/source/include/rocprofiler-sdk/cxx/perfetto.hpp new file mode 100644 index 0000000000..dabc494eb5 --- /dev/null +++ b/source/include/rocprofiler-sdk/cxx/perfetto.hpp @@ -0,0 +1,168 @@ +// MIT License +// +// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + +#pragma once + +#include + +#include +#include +#include +#include +#include +#include +#include + +#define ROCPROFILER_DEFINE_PERFETTO_CATEGORY(NAME, DESC, ...) \ + namespace rocprofiler \ + { \ + namespace sdk \ + { \ + template <> \ + struct perfetto_category<__VA_ARGS__> \ + { \ + static constexpr auto name = NAME; \ + static constexpr auto description = DESC; \ + }; \ + } \ + } + +#define ROCPROFILER_DEFINE_CATEGORY(NS, VALUE, DESC) \ + namespace rocprofiler \ + { \ + namespace sdk \ + { \ + namespace NS \ + { \ + struct VALUE; \ + } \ + } \ + } \ + ROCPROFILER_DEFINE_PERFETTO_CATEGORY(#VALUE, DESC, NS::VALUE) + +#define ROCPROFILER_PERFETTO_CATEGORY(TYPE) \ + ::perfetto::Category(::rocprofiler::sdk::perfetto_category<::rocprofiler::sdk::TYPE>::name) \ + .SetDescription( \ + ::rocprofiler::sdk::perfetto_category<::rocprofiler::sdk::TYPE>::description) + +namespace rocprofiler +{ +namespace sdk +{ +template +struct perfetto_category; +} // namespace sdk +} // namespace rocprofiler + +ROCPROFILER_DEFINE_CATEGORY(category, hsa_api, "HSA API function") +ROCPROFILER_DEFINE_CATEGORY(category, hip_api, "HIP API function") +ROCPROFILER_DEFINE_CATEGORY(category, marker_api, "Marker API region") +ROCPROFILER_DEFINE_CATEGORY(category, kernel_dispatch, "GPU kernel dispatch") +ROCPROFILER_DEFINE_CATEGORY(category, memory_copy, "Async memory copy") + +#define ROCPROFILER_PERFETTO_CATEGORIES \ + ROCPROFILER_PERFETTO_CATEGORY(category::hsa_api), \ + ROCPROFILER_PERFETTO_CATEGORY(category::hip_api), \ + ROCPROFILER_PERFETTO_CATEGORY(category::marker_api), \ + ROCPROFILER_PERFETTO_CATEGORY(category::kernel_dispatch), \ + ROCPROFILER_PERFETTO_CATEGORY(category::memory_copy) + +#include + +PERFETTO_DEFINE_CATEGORIES(ROCPROFILER_PERFETTO_CATEGORIES); + +namespace rocprofiler +{ +namespace sdk +{ +using perfetto_event_context_t = ::perfetto::EventContext; + +template +auto +add_perfetto_annotation(perfetto_event_context_t& ctx, Np&& _name, Tp&& _val) +{ + namespace mpl = ::rocprofiler::sdk::mpl; + + using named_type = mpl::unqualified_identity_t; + using value_type = mpl::unqualified_identity_t; + + static_assert(mpl::is_string_type::value, "Error! name is not a string type"); + + auto _get_dbg = [&]() { + auto* _dbg = ctx.event()->add_debug_annotations(); + _dbg->set_name(std::string_view{std::forward(_name)}.data()); + return _dbg; + }; + + if constexpr(std::is_same::value) + { + _get_dbg()->set_string_value(_val.data()); + } + else if constexpr(mpl::is_string_type::value) + { + _get_dbg()->set_string_value(std::forward(_val)); + } + else if constexpr(std::is_same::value) + { + _get_dbg()->set_bool_value(_val); + } + else if constexpr(std::is_enum::value) + { + _get_dbg()->set_int_value(static_cast(_val)); + } + else if constexpr(std::is_floating_point::value) + { + _get_dbg()->set_double_value(static_cast(_val)); + } + else if constexpr(std::is_integral::value) + { + if constexpr(std::is_unsigned::value) + { + _get_dbg()->set_uint_value(_val); + } + else + { + _get_dbg()->set_int_value(_val); + } + } + else if constexpr(std::is_pointer::value) + { + _get_dbg()->set_pointer_value(reinterpret_cast(_val)); + } + else if constexpr(mpl::can_stringify::value) + { + auto _ss = std::stringstream{}; + _ss << std::forward(_val); + _get_dbg()->set_string_value(_ss.str()); + } + else + { + static_assert(std::is_empty::value, "Error! unsupported data type"); + } +} +} // namespace sdk +} // namespace rocprofiler + +#undef ROCPROFILER_DEFINE_PERFETTO_CATEGORY +#undef ROCPROFILER_DEFINE_CATEGORY +#undef ROCPROFILER_PERFETTO_CATEGORY +#undef ROCPROFILER_PERFETTO_CATEGORIES diff --git a/source/lib/rocprofiler-sdk-tool/CMakeLists.txt b/source/lib/rocprofiler-sdk-tool/CMakeLists.txt index ed2c13621b..fe17d41b32 100644 --- a/source/lib/rocprofiler-sdk-tool/CMakeLists.txt +++ b/source/lib/rocprofiler-sdk-tool/CMakeLists.txt @@ -11,6 +11,7 @@ set(TOOL_HEADERS domain_type.hpp generateCSV.hpp generateJSON.hpp + generatePerfetto.hpp helper.hpp output_file.hpp statistics.hpp @@ -22,6 +23,7 @@ set(TOOL_SOURCES domain_type.cpp generateCSV.cpp generateJSON.cpp + generatePerfetto.cpp helper.cpp main.c output_file.cpp @@ -36,9 +38,13 @@ add_subdirectory(plugins) target_link_libraries( rocprofiler-sdk-tool - PRIVATE rocprofiler::rocprofiler-shared-library rocprofiler::rocprofiler-headers - rocprofiler::rocprofiler-build-flags rocprofiler::rocprofiler-memcheck - rocprofiler::rocprofiler-common-library rocprofiler::rocprofiler-cereal) + PRIVATE rocprofiler::rocprofiler-shared-library + rocprofiler::rocprofiler-headers + rocprofiler::rocprofiler-build-flags + rocprofiler::rocprofiler-memcheck + rocprofiler::rocprofiler-common-library + rocprofiler::rocprofiler-cereal + rocprofiler::rocprofiler-perfetto) set_target_properties( rocprofiler-sdk-tool diff --git a/source/lib/rocprofiler-sdk-tool/config.cpp b/source/lib/rocprofiler-sdk-tool/config.cpp index 367e783e6d..e7ffe360af 100644 --- a/source/lib/rocprofiler-sdk-tool/config.cpp +++ b/source/lib/rocprofiler-sdk-tool/config.cpp @@ -271,10 +271,11 @@ config::config() break; } - csv_output = entries.count("CSV") > 0 || entries.empty(); - json_output = entries.count("JSON") > 0; + csv_output = entries.count("CSV") > 0 || entries.empty(); + json_output = entries.count("JSON") > 0; + pftrace_output = entries.count("PFTRACE") > 0; - const auto supported_formats = std::set{"CSV", "JSON"}; + const auto supported_formats = std::set{"CSV", "JSON", "PFTRACE"}; for(const auto& itr : entries) { LOG_IF(FATAL, supported_formats.count(itr) == 0) diff --git a/source/lib/rocprofiler-sdk-tool/config.hpp b/source/lib/rocprofiler-sdk-tool/config.hpp index dc5cca7d1d..2c2cae451b 100644 --- a/source/lib/rocprofiler-sdk-tool/config.hpp +++ b/source/lib/rocprofiler-sdk-tool/config.hpp @@ -75,6 +75,7 @@ struct config bool stats = get_env("ROCPROF_STATS", false); bool csv_output = false; bool json_output = false; + bool pftrace_output = false; int mpi_size = get_mpi_size(); int mpi_rank = get_mpi_rank(); std::string output_path = get_env("ROCPROF_OUTPUT_PATH", fs::current_path().string()); diff --git a/source/lib/rocprofiler-sdk-tool/generateCSV.cpp b/source/lib/rocprofiler-sdk-tool/generateCSV.cpp index cd8da07a8a..ca72d0075a 100644 --- a/source/lib/rocprofiler-sdk-tool/generateCSV.cpp +++ b/source/lib/rocprofiler-sdk-tool/generateCSV.cpp @@ -28,6 +28,7 @@ #include #include +#include #include #include diff --git a/source/lib/rocprofiler-sdk-tool/generatePerfetto.cpp b/source/lib/rocprofiler-sdk-tool/generatePerfetto.cpp new file mode 100644 index 0000000000..73725fe79b --- /dev/null +++ b/source/lib/rocprofiler-sdk-tool/generatePerfetto.cpp @@ -0,0 +1,529 @@ +// MIT License +// +// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + +#include "generatePerfetto.hpp" +#include "helper.hpp" +#include "output_file.hpp" + +#include "lib/common/utility.hpp" + +#include +#include +#include +#include +#include +#include +#include + +#include +#include +#include +#include + +namespace rocprofiler +{ +namespace tool +{ +namespace +{ +auto main_tid = common::get_tid(); + +template +size_t +get_hash_id(Tp&& _val) +{ + if constexpr(!std::is_pointer::value) + return std::hash{}(std::forward(_val)); + else if constexpr(std::is_same::value) + return get_hash_id(std::string_view{_val}); + else + return get_hash_id(*_val); +} +} // namespace + +void +write_perfetto( + tool_table* tool_functions, + uint64_t /*pid*/, + std::vector agent_data, + std::deque* hip_api_data, + std::deque* hsa_api_data, + std::deque* kernel_dispatch_data, + std::deque* memory_copy_data, + std::deque* marker_api_data, + std::deque* /*scratch_api_data*/) +{ + namespace sdk = ::rocprofiler::sdk; + + auto agents_map = std::unordered_map{}; + for(auto itr : agent_data) + agents_map.emplace(itr.id, itr); + + auto args = ::perfetto::TracingInitArgs{}; + auto track_event_cfg = ::perfetto::protos::gen::TrackEventConfig{}; + auto cfg = ::perfetto::TraceConfig{}; + + // environment settings + auto shmem_size_hint = size_t{64}; + auto buffer_size_kb = size_t{1024000}; + + auto* buffer_config = cfg.add_buffers(); + buffer_config->set_size_kb(buffer_size_kb); + buffer_config->set_fill_policy( + ::perfetto::protos::gen::TraceConfig_BufferConfig_FillPolicy_DISCARD); + + auto* ds_cfg = cfg.add_data_sources()->mutable_config(); + ds_cfg->set_name("track_event"); // this MUST be track_event + ds_cfg->set_track_event_config_raw(track_event_cfg.SerializeAsString()); + + args.shmem_size_hint_kb = shmem_size_hint; + args.backends |= ::perfetto::kInProcessBackend; + + ::perfetto::Tracing::Initialize(args); + ::perfetto::TrackEvent::Register(); + + auto tracing_session = ::perfetto::Tracing::NewTrace(); + + tracing_session->Setup(cfg); + tracing_session->StartBlocking(); + + auto tids = std::set{}; + auto demangled = std::unordered_map{}; + auto agent_thread_ids = std::unordered_map>{}; + auto agent_queue_ids = + std::unordered_map>{}; + auto thread_indexes = std::unordered_map{}; + auto kernel_sym_data = get_kernel_symbol_data(); + + auto thread_tracks = std::unordered_map{}; + auto agent_thread_tracks = + std::unordered_map>{}; + auto agent_queue_tracks = + std::unordered_map>{}; + + auto _get_agent = [&agent_data](rocprofiler_agent_id_t _id) -> const rocprofiler_agent_t* { + for(const auto& itr : agent_data) + { + if(_id == itr.id) return &itr; + } + return CHECK_NOTNULL(nullptr); + }; + + { + for(auto itr : *hsa_api_data) + tids.emplace(itr.thread_id); + for(auto itr : *hip_api_data) + tids.emplace(itr.thread_id); + for(auto itr : *marker_api_data) + tids.emplace(itr.thread_id); + + for(auto itr : *memory_copy_data) + agent_thread_ids[itr.dst_agent_id].emplace(itr.thread_id); + + for(auto itr : *kernel_dispatch_data) + agent_queue_ids[itr.dispatch_info.agent_id].emplace(itr.dispatch_info.queue_id); + } + + uint64_t nthrn = 0; + for(auto itr : tids) + { + if(itr == main_tid) + { + thread_indexes.emplace(main_tid, 0); + thread_tracks.emplace(main_tid, ::perfetto::ThreadTrack::Current()); + } + else + { + auto _idx = ++nthrn; + thread_indexes.emplace(itr, _idx); + auto _track = ::perfetto::Track{itr}; + auto _desc = _track.Serialize(); + auto _namess = std::stringstream{}; + _namess << "THREAD " << _idx << " (" << itr << ")"; + _desc.set_name(_namess.str()); + perfetto::TrackEvent::SetTrackDescriptor(_track, _desc); + + thread_tracks.emplace(itr, _track); + } + } + + for(const auto& itr : agent_thread_ids) + { + const auto* _agent = _get_agent(itr.first); + + for(auto titr : itr.second) + { + auto _namess = std::stringstream{}; + _namess << "COPY to [" << _agent->logical_node_id << "] THREAD [" + << thread_indexes.at(titr) << "]"; + + if(_agent->type == ROCPROFILER_AGENT_TYPE_CPU) + _namess << " CPU"; + else if(_agent->type == ROCPROFILER_AGENT_TYPE_GPU) + _namess << " GPU"; + + auto _track = ::perfetto::Track{get_hash_id(_namess.str())}; + auto _desc = _track.Serialize(); + _desc.set_name(_namess.str()); + + perfetto::TrackEvent::SetTrackDescriptor(_track, _desc); + + agent_thread_tracks[itr.first].emplace(titr, _track); + } + } + + for(const auto& aitr : agent_queue_ids) + { + uint32_t nqueue = 0; + for(auto qitr : aitr.second) + { + const auto* _agent = _get_agent(aitr.first); + + auto _namess = std::stringstream{}; + _namess << "COMPUTE [" << _agent->logical_node_id << "] QUEUE [" << nqueue++ << "] "; + + if(_agent->type == ROCPROFILER_AGENT_TYPE_CPU) + _namess << "CPU"; + else if(_agent->type == ROCPROFILER_AGENT_TYPE_GPU) + _namess << "GPU"; + + auto _track = ::perfetto::Track{get_hash_id(_namess.str())}; + auto _desc = _track.Serialize(); + _desc.set_name(_namess.str()); + + perfetto::TrackEvent::SetTrackDescriptor(_track, _desc); + + agent_queue_tracks[aitr.first].emplace(qitr, _track); + } + } + + // trace events + { + auto buffer_names = sdk::get_buffer_tracing_names(); + auto callbk_name_info = sdk::get_callback_tracing_names(); + + for(auto itr : *hsa_api_data) + { + auto name = buffer_names.at(itr.kind, itr.operation); + auto& track = thread_tracks.at(itr.thread_id); + + TRACE_EVENT_BEGIN(sdk::perfetto_category::name, + ::perfetto::StaticString(name.data()), + track, + itr.start_timestamp, + ::perfetto::Flow::ProcessScoped(itr.correlation_id.internal), + "begin_ns", + itr.start_timestamp, + "end_ns", + itr.end_timestamp, + "delta_ns", + (itr.end_timestamp - itr.start_timestamp), + "tid", + itr.thread_id, + "kind", + itr.kind, + "operation", + itr.operation, + "corr_id", + itr.correlation_id.internal); + TRACE_EVENT_END( + sdk::perfetto_category::name, track, itr.end_timestamp); + } + + for(auto itr : *hip_api_data) + { + auto name = buffer_names.at(itr.kind, itr.operation); + auto& track = thread_tracks.at(itr.thread_id); + + TRACE_EVENT_BEGIN(sdk::perfetto_category::name, + ::perfetto::StaticString(name.data()), + track, + itr.start_timestamp, + ::perfetto::Flow::ProcessScoped(itr.correlation_id.internal), + "begin_ns", + itr.start_timestamp, + "end_ns", + itr.end_timestamp, + "delta_ns", + (itr.end_timestamp - itr.start_timestamp), + "tid", + itr.thread_id, + "kind", + itr.kind, + "operation", + itr.operation, + "corr_id", + itr.correlation_id.internal); + TRACE_EVENT_END( + sdk::perfetto_category::name, track, itr.end_timestamp); + } + + for(auto itr : *marker_api_data) + { + auto& track = thread_tracks.at(itr.thread_id); + auto name = (itr.kind == ROCPROFILER_BUFFER_TRACING_MARKER_CORE_API && + itr.operation != ROCPROFILER_MARKER_CORE_API_ID_roctxGetThreadId) + ? tool_functions->tool_get_roctx_msg_fn(itr.correlation_id.internal) + : buffer_names.at(itr.kind, itr.operation); + + TRACE_EVENT_BEGIN(sdk::perfetto_category::name, + ::perfetto::StaticString(name.data()), + track, + itr.start_timestamp, + ::perfetto::Flow::ProcessScoped(itr.correlation_id.internal), + "begin_ns", + itr.start_timestamp, + "end_ns", + itr.end_timestamp, + "delta_ns", + (itr.end_timestamp - itr.start_timestamp), + "tid", + itr.thread_id, + "kind", + itr.kind, + "operation", + itr.operation, + "corr_id", + itr.correlation_id.internal); + TRACE_EVENT_END( + sdk::perfetto_category::name, track, itr.end_timestamp); + } + + for(auto itr : *memory_copy_data) + { + auto name = buffer_names.at(itr.kind, itr.operation); + auto& track = agent_thread_tracks.at(itr.dst_agent_id).at(itr.thread_id); + + TRACE_EVENT_BEGIN(sdk::perfetto_category::name, + ::perfetto::StaticString(name.data()), + track, + itr.start_timestamp, + ::perfetto::Flow::ProcessScoped(itr.correlation_id.internal), + "begin_ns", + itr.start_timestamp, + "end_ns", + itr.end_timestamp, + "delta_ns", + (itr.end_timestamp - itr.start_timestamp), + "kind", + itr.kind, + "operation", + itr.operation, + "src_agent", + agents_map.at(itr.src_agent_id).logical_node_id, + "dst_agent", + agents_map.at(itr.dst_agent_id).logical_node_id, + "copy_bytes", + itr.bytes, + "corr_id", + itr.correlation_id.internal, + "tid", + itr.thread_id); + TRACE_EVENT_END( + sdk::perfetto_category::name, track, itr.end_timestamp); + } + + for(auto itr : *kernel_dispatch_data) + { + const auto& info = itr.dispatch_info; + const kernel_symbol_data* sym = nullptr; + for(const auto& kitr : kernel_sym_data) + { + if(kitr.kernel_id == info.kernel_id) + { + sym = &kitr; + break; + } + } + + CHECK(sym != nullptr); + + auto name = std::string_view{sym->kernel_name}; + auto& track = agent_queue_tracks.at(info.agent_id).at(info.queue_id); + + if(demangled.find(name) == demangled.end()) + { + demangled.emplace(name, common::cxx_demangle(name)); + } + + TRACE_EVENT_BEGIN(sdk::perfetto_category::name, + ::perfetto::StaticString(demangled.at(name).c_str()), + track, + itr.start_timestamp, + ::perfetto::Flow::ProcessScoped(itr.correlation_id.internal), + "begin_ns", + itr.start_timestamp, + "end_ns", + itr.end_timestamp, + "delta_ns", + (itr.end_timestamp - itr.start_timestamp), + "kind", + itr.kind, + "agent", + agents_map.at(info.agent_id).logical_node_id, + "corr_id", + itr.correlation_id.internal, + "queue", + info.queue_id.handle, + "tid", + itr.thread_id, + "kernel_id", + info.kernel_id, + "private_segment_size", + info.private_segment_size, + "group_segment_size", + info.group_segment_size, + "workgroup_size", + info.workgroup_size.x * info.workgroup_size.y * info.workgroup_size.z, + "grid_size", + info.grid_size.x * info.grid_size.y * info.grid_size.z); + TRACE_EVENT_END(sdk::perfetto_category::name, + track, + itr.end_timestamp); + } + } + + // counter tracks + { + // memory copy counter track + auto mem_cpy_endpoints = std::map>{}; + auto mem_cpy_extremes = std::pair{}; + for(auto itr : *memory_copy_data) + { + uint64_t _mean_timestamp = + itr.start_timestamp + (0.5 * (itr.end_timestamp - itr.start_timestamp)); + + mem_cpy_endpoints[itr.dst_agent_id].emplace(itr.start_timestamp - 1000, 0); + mem_cpy_endpoints[itr.dst_agent_id].emplace(itr.start_timestamp, 0); + mem_cpy_endpoints[itr.dst_agent_id].emplace(_mean_timestamp, 0); + mem_cpy_endpoints[itr.dst_agent_id].emplace(itr.end_timestamp, 0); + mem_cpy_endpoints[itr.dst_agent_id].emplace(itr.end_timestamp + 1000, 0); + + mem_cpy_extremes = std::make_pair(std::min(mem_cpy_extremes.first, itr.start_timestamp), + std::max(mem_cpy_extremes.second, itr.end_timestamp)); + } + + for(auto itr : *memory_copy_data) + { + auto mbeg = mem_cpy_endpoints.at(itr.dst_agent_id).lower_bound(itr.start_timestamp); + auto mend = mem_cpy_endpoints.at(itr.dst_agent_id).upper_bound(itr.end_timestamp); + + LOG_IF(FATAL, mbeg == mend) << "Missing range for timestamp [" << itr.start_timestamp + << ", " << itr.end_timestamp << "]"; + + for(auto mitr = mbeg; mitr != mend; ++mitr) + mitr->second += itr.bytes; + } + + constexpr auto bytes_multiplier = 1024; + + auto mem_cpy_tracks = + std::unordered_map{}; + auto mem_cpy_cnt_names = std::vector{}; + mem_cpy_cnt_names.reserve(mem_cpy_endpoints.size()); + for(auto& mitr : mem_cpy_endpoints) + { + mem_cpy_endpoints[mitr.first].emplace(mem_cpy_extremes.first - 5000, 0); + mem_cpy_endpoints[mitr.first].emplace(mem_cpy_extremes.second + 5000, 0); + + auto _track_name = std::stringstream{}; + const auto* _agent = _get_agent(mitr.first); + + if(_agent->type == ROCPROFILER_AGENT_TYPE_CPU) + _track_name << "COPY BYTES to [" << _agent->logical_node_id << "] CPU"; + else if(_agent->type == ROCPROFILER_AGENT_TYPE_GPU) + _track_name << "COPY BYTES to [" << _agent->logical_node_id << "] GPU"; + + constexpr auto _unit = ::perfetto::CounterTrack::Unit::UNIT_SIZE_BYTES; + auto& _name = mem_cpy_cnt_names.emplace_back(_track_name.str()); + mem_cpy_tracks.emplace(mitr.first, + ::perfetto::CounterTrack{_name.c_str()} + .set_unit(_unit) + .set_unit_multiplier(bytes_multiplier) + .set_is_incremental(false)); + } + + for(auto& mitr : mem_cpy_endpoints) + { + for(auto itr : mitr.second) + { + TRACE_COUNTER(sdk::perfetto_category::name, + mem_cpy_tracks.at(mitr.first), + itr.first, + itr.second / bytes_multiplier); + } + } + } + + ::perfetto::TrackEvent::Flush(); + tracing_session->FlushBlocking(); + tracing_session->StopBlocking(); + + auto filename = std::string{"results"}; + auto cleanup = std::function{}; + std::ostream* ofs = nullptr; + + std::tie(ofs, cleanup) = get_output_stream(filename, ".pftrace"); + + auto amount_read = std::atomic{0}; + auto is_done = std::promise{}; + auto _mtx = std::mutex{}; + auto _reader = [ofs, &_mtx, &is_done, &amount_read]( + ::perfetto::TracingSession::ReadTraceCallbackArgs _args) { + auto _lk = std::unique_lock{_mtx}; + if(_args.data && _args.size > 0) + { + ROCP_TRACE << "Writing " << _args.size << " B to trace..."; + // Write the trace data into file + ofs->write(_args.data, _args.size); + amount_read += _args.size; + } + ROCP_INFO_IF(!_args.has_more && amount_read > 0) + << "Wrote " << amount_read << " B to perfetto trace file"; + if(!_args.has_more) is_done.set_value(); + }; + + for(size_t i = 0; i < 2; ++i) + { + ROCP_TRACE << "Reading trace..."; + amount_read = 0; + is_done = std::promise{}; + tracing_session->ReadTrace(_reader); + is_done.get_future().wait(); + } + + ROCP_TRACE << "Destroying tracing session..."; + tracing_session.reset(); + + ROCP_TRACE << "Flushing trace output stream..."; + (*ofs) << std::flush; + + ROCP_TRACE << "Destroying trace output stream..."; + if(cleanup) cleanup(ofs); +} + +} // namespace tool +} // namespace rocprofiler + +PERFETTO_TRACK_EVENT_STATIC_STORAGE(); diff --git a/source/lib/rocprofiler-sdk-tool/generatePerfetto.hpp b/source/lib/rocprofiler-sdk-tool/generatePerfetto.hpp new file mode 100644 index 0000000000..80837931bd --- /dev/null +++ b/source/lib/rocprofiler-sdk-tool/generatePerfetto.hpp @@ -0,0 +1,45 @@ +// MIT License +// +// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + +#pragma once + +#include "helper.hpp" + +#include + +namespace rocprofiler +{ +namespace tool +{ +void +write_perfetto( + tool_table* tool_functions, + uint64_t pid, + std::vector agent_data, + std::deque* hip_api_data, + std::deque* hsa_api_data, + std::deque* kernel_dispatch_data, + std::deque* memory_copy_data, + std::deque* marker_api_data, + std::deque* scratch_api_data); +} // namespace tool +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk-tool/tool.cpp b/source/lib/rocprofiler-sdk-tool/tool.cpp index 807023c18c..ebec4268f6 100644 --- a/source/lib/rocprofiler-sdk-tool/tool.cpp +++ b/source/lib/rocprofiler-sdk-tool/tool.cpp @@ -26,6 +26,7 @@ #include "domain_type.hpp" #include "generateCSV.hpp" #include "generateJSON.hpp" +#include "generatePerfetto.hpp" #include "helper.hpp" #include "output_file.hpp" #include "tmp_file.hpp" @@ -1433,6 +1434,19 @@ tool_fini(void* /*tool_data*/) &scratch_memory_output.element_data); } + if(tool::get_config().pftrace_output) + { + rocprofiler::tool::write_perfetto(tool_functions, + getpid(), + _agents, + &hip_output.element_data, + &hsa_output.element_data, + &kernel_dispatch_output.element_data, + &memory_copy_output.element_data, + &marker_output.element_data, + &scratch_memory_output.element_data); + } + auto destroy_output = [](auto& _buffered_output_v) { _buffered_output_v.destroy(); }; destroy_output(kernel_dispatch_output); diff --git a/tests/bin/hip-in-libraries/hip-in-libraries.cpp b/tests/bin/hip-in-libraries/hip-in-libraries.cpp index 9e78ec7692..24d8cc0f16 100644 --- a/tests/bin/hip-in-libraries/hip-in-libraries.cpp +++ b/tests/bin/hip-in-libraries/hip-in-libraries.cpp @@ -56,7 +56,7 @@ namespace { using auto_lock_t = std::unique_lock; auto print_lock = std::mutex{}; -size_t nqueues = 8; +size_t nqueues = 4; size_t nthreads = 4; size_t nitr = 500; size_t nsync = 10; @@ -112,7 +112,8 @@ main(int argc, char** argv) { auto vector_ops_thread = std::thread{run_vector_ops, nthreads, nqueues}; - auto transpose_thread = std::thread{run_transpose, nthreads, nitr, nsync}; + std::this_thread::sleep_for(std::chrono::milliseconds{100}); + auto transpose_thread = std::thread{run_transpose, nthreads, nitr, nsync}; vector_ops_thread.join(); transpose_thread.join(); diff --git a/tests/common/CMakeLists.txt b/tests/common/CMakeLists.txt index 890bc3eb1e..db04da9670 100644 --- a/tests/common/CMakeLists.txt +++ b/tests/common/CMakeLists.txt @@ -84,32 +84,40 @@ if(NOT TARGET rocprofiler::rocprofiler-cereal) endif() endif() -# perfetto -fetchcontent_declare( - perfetto - GIT_REPOSITORY https://android.googlesource.com/platform/external/perfetto - GIT_TAG v42.0 - SOURCE_DIR ${PROJECT_BINARY_DIR}/external/perfetto BINARY_DIR - ${PROJECT_BINARY_DIR}/external/build/perfetto-build SUBBUILD_DIR - ${PROJECT_BINARY_DIR}/external/build/perfetto-subdir) +if(NOT TARGET rocprofiler::rocprofiler-perfetto) + # perfetto + fetchcontent_declare( + perfetto + GIT_REPOSITORY https://android.googlesource.com/platform/external/perfetto + GIT_TAG v44.0 + SOURCE_DIR ${PROJECT_BINARY_DIR}/external/perfetto BINARY_DIR + ${PROJECT_BINARY_DIR}/external/build/perfetto-build SUBBUILD_DIR + ${PROJECT_BINARY_DIR}/external/build/perfetto-subdir) -fetchcontent_getproperties(perfetto) + fetchcontent_getproperties(perfetto) -if(NOT perfetto_POPULATED) - fetchcontent_populate(perfetto) + if(NOT perfetto_POPULATED) + fetchcontent_populate(perfetto) + endif() + + add_library(rocprofiler-tests-perfetto STATIC) + add_library(rocprofiler::tests-perfetto ALIAS rocprofiler-tests-perfetto) + target_sources( + rocprofiler-tests-perfetto + PRIVATE ${PROJECT_BINARY_DIR}/external/perfetto/sdk/perfetto.h + ${PROJECT_BINARY_DIR}/external/perfetto/sdk/perfetto.cc) + target_include_directories( + rocprofiler-tests-perfetto SYSTEM + INTERFACE $) + set_target_properties(rocprofiler-tests-perfetto PROPERTIES POSITION_INDEPENDENT_CODE + ON) +else() + add_library(rocprofiler-tests-perfetto INTERFACE) + add_library(rocprofiler::tests-perfetto ALIAS rocprofiler-tests-perfetto) + target_link_libraries(rocprofiler-tests-perfetto + INTERFACE rocprofiler::rocprofiler-perfetto) endif() -add_library(rocprofiler-tests-perfetto STATIC) -add_library(rocprofiler::tests-perfetto ALIAS rocprofiler-tests-perfetto) -target_sources( - rocprofiler-tests-perfetto - PRIVATE ${PROJECT_BINARY_DIR}/external/perfetto/sdk/perfetto.h - ${PROJECT_BINARY_DIR}/external/perfetto/sdk/perfetto.cc) -target_include_directories( - rocprofiler-tests-perfetto SYSTEM - INTERFACE $) -set_target_properties(rocprofiler-tests-perfetto PROPERTIES POSITION_INDEPENDENT_CODE ON) - # common utilities cmake_path(GET CMAKE_CURRENT_SOURCE_DIR PARENT_PATH COMMON_LIBRARY_INCLUDE_DIR) diff --git a/tests/common/perfetto.hpp b/tests/common/perfetto.hpp index f8fe8e92c4..f6d9084aac 100644 --- a/tests/common/perfetto.hpp +++ b/tests/common/perfetto.hpp @@ -22,204 +22,4 @@ #pragma once -#include -#include -#include -#include -#include -#include -#include - -#define ROCPROFILER_DEFINE_NAME_TRAIT(NAME, DESC, ...) \ - namespace rocprofiler \ - { \ - template <> \ - struct perfetto_category<__VA_ARGS__> \ - { \ - static constexpr auto value = NAME; \ - static constexpr auto description = DESC; \ - }; \ - } - -namespace rocprofiler -{ -template -struct perfetto_category; - -namespace trait -{ -template -using name = perfetto_category; -} -} // namespace rocprofiler - -#define ROCPROFILER_DEFINE_NS_API(NS, NAME) \ - namespace rocprofiler \ - { \ - namespace NS \ - { \ - struct NAME; \ - } \ - } - -#define ROCPROFILER_DEFINE_CATEGORY(NS, VALUE, DESC) \ - ROCPROFILER_DEFINE_NS_API(NS, VALUE) \ - ROCPROFILER_DEFINE_NAME_TRAIT(#VALUE, DESC, NS::VALUE) - -ROCPROFILER_DEFINE_CATEGORY(category, hsa_api, "HSA API function") -ROCPROFILER_DEFINE_CATEGORY(category, hip_api, "HIP API function") -ROCPROFILER_DEFINE_CATEGORY(category, marker_api, "Marker API region") -ROCPROFILER_DEFINE_CATEGORY(category, kernel_dispatch, "GPU kernel dispatch") -ROCPROFILER_DEFINE_CATEGORY(category, memory_copy, "Async memory copy") - -#define ROCPROFILER_PERFETTO_CATEGORY(TYPE) \ - ::perfetto::Category(::rocprofiler::perfetto_category<::rocprofiler::TYPE>::value) \ - .SetDescription(::rocprofiler::perfetto_category<::rocprofiler::TYPE>::description) - -#define ROCPROFILER_PERFETTO_CATEGORIES \ - ROCPROFILER_PERFETTO_CATEGORY(category::hsa_api), \ - ROCPROFILER_PERFETTO_CATEGORY(category::hip_api), \ - ROCPROFILER_PERFETTO_CATEGORY(category::marker_api), \ - ROCPROFILER_PERFETTO_CATEGORY(category::kernel_dispatch), \ - ROCPROFILER_PERFETTO_CATEGORY(category::memory_copy) - -#include - -PERFETTO_DEFINE_CATEGORIES(ROCPROFILER_PERFETTO_CATEGORIES); - -namespace concepts -{ -template -struct is_string_type : std::false_type -{}; - -template <> -struct is_string_type : std::true_type -{}; - -template <> -struct is_string_type : std::true_type -{}; - -template <> -struct is_string_type : std::true_type -{}; - -template <> -struct is_string_type : std::true_type -{}; - -template -struct is_string_type : is_string_type> -{}; - -template -struct is_string_type : is_string_type> -{}; - -template -struct is_string_type : is_string_type> -{}; - -template -struct is_string_type : is_string_type> -{}; - -template -struct is_string_type : is_string_type> -{}; - -template -struct is_string_type : is_string_type> -{}; - -template -struct unqualified_type -{ - using type = std::remove_reference_t>>; -}; - -template -using unqualified_type_t = typename unqualified_type::type; - -template -struct can_stringify -{ -private: - static constexpr auto sfinae(int) - -> decltype(std::declval() << std::declval(), bool()) - { - return true; - } - - static constexpr auto sfinae(long) { return false; } - -public: - static constexpr bool value = sfinae(0); - constexpr auto operator()() const { return sfinae(0); } -}; -} // namespace concepts - -using perfetto_event_context_t = ::perfetto::EventContext; - -template -auto -add_perfetto_annotation(perfetto_event_context_t& ctx, Np&& _name, Tp&& _val) -{ - using named_type = concepts::unqualified_type_t; - using value_type = concepts::unqualified_type_t; - - static_assert(concepts::is_string_type::value, "Error! name is not a string type"); - - auto _get_dbg = [&]() { - auto* _dbg = ctx.event()->add_debug_annotations(); - _dbg->set_name(std::string_view{std::forward(_name)}.data()); - return _dbg; - }; - - if constexpr(std::is_same::value) - { - _get_dbg()->set_string_value(_val.data()); - } - else if constexpr(concepts::is_string_type::value) - { - _get_dbg()->set_string_value(std::forward(_val)); - } - else if constexpr(std::is_same::value) - { - _get_dbg()->set_bool_value(_val); - } - else if constexpr(std::is_enum::value) - { - _get_dbg()->set_int_value(static_cast(_val)); - } - else if constexpr(std::is_floating_point::value) - { - _get_dbg()->set_double_value(static_cast(_val)); - } - else if constexpr(std::is_integral::value) - { - if constexpr(std::is_unsigned::value) - { - _get_dbg()->set_uint_value(_val); - } - else - { - _get_dbg()->set_int_value(_val); - } - } - else if constexpr(std::is_pointer::value) - { - _get_dbg()->set_pointer_value(reinterpret_cast(_val)); - } - else if constexpr(concepts::can_stringify::value) - { - auto _ss = std::stringstream{}; - _ss << std::forward(_val); - _get_dbg()->set_string_value(_ss.str()); - } - else - { - static_assert(std::is_empty::value, "Error! unsupported data type"); - } -} +#include diff --git a/tests/pytest-packages/CMakeLists.txt b/tests/pytest-packages/CMakeLists.txt index f5c11841b7..fc347a10cd 100644 --- a/tests/pytest-packages/CMakeLists.txt +++ b/tests/pytest-packages/CMakeLists.txt @@ -4,13 +4,20 @@ set(PACKAGE_OUTPUT_DIR ${ROCPROFILER_SDK_TESTS_BINARY_DIR}/pytest-packages/rocprofiler_sdk) -set(PYTEST_UTILS_PYTHON_SOURCES __init__.py dotdict.py) + +set(PYTEST_UTILS_PYTHON_SOURCES __init__.py dotdict.py perfetto_reader.py) +set(TESTS_PYTHON_SOURCES __init__.py rocprofv3.py) foreach(_FILE ${PYTEST_UTILS_PYTHON_SOURCES}) configure_file(${CMAKE_CURRENT_LIST_DIR}/pytest_utils/${_FILE} ${PACKAGE_OUTPUT_DIR}/pytest_utils/${_FILE} COPYONLY) endforeach() +foreach(_FILE ${TESTS_PYTHON_SOURCES}) + configure_file(${CMAKE_CURRENT_LIST_DIR}/tests/${_FILE} + ${PACKAGE_OUTPUT_DIR}/tests/${_FILE} COPYONLY) +endforeach() + file( WRITE "${PACKAGE_OUTPUT_DIR}/__init__.py" "# diff --git a/tests/pytest-packages/pytest_utils/perfetto_reader.py b/tests/pytest-packages/pytest_utils/perfetto_reader.py new file mode 100644 index 0000000000..a1f4d820d1 --- /dev/null +++ b/tests/pytest-packages/pytest_utils/perfetto_reader.py @@ -0,0 +1,750 @@ +# MIT License +# +# Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to deal +# in the Software without restriction, including without limitation the rights +# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +# copies of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in all +# copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +# SOFTWARE. + +from __future__ import absolute_import + +import re +import pandas as pd + +from collections import OrderedDict +from perfetto.trace_processor import TraceProcessor + + +class PerfettoReader: + """Read in perfetto protobuf output""" + + def __init__(self, filename, select=None, **kwargs): + """Arguments: + filename (str or list/tuple of str): + Valid arguments should be a list of files + + verbose (int): + Information about processes, threads, categories, performance, etc. + + report (list of str): + alternative for verbose. Accepts: + + - category (report categories in files) + - process (report process info in files) + - threads (report thread info in files) + - track_ids (report track id info in files) + - profile (report timing info from processing) + + exclude_category (list of str): + Every slice has an associated category, e.g., slices from sampling may be + in "sampling" category, slices from instrumentation may be in "instrumentation" + category. The categories in protobuf is specific to the tool that generated the + protobuf. Use this option to exclude the slices from specific categories + + include_category (list of str): + Every slice has an associated category, e.g., slices from sampling may be + in "sampling" category, slices from instrumentation may be in "instrumentation" + category. The categories in protobuf is specific to the tool that generated the + protobuf. Use this option to restrict the slices to the specified categories + + default_categories (list of str): + Use this option as a safety value when include/exclude are used. These categories + are used when include/exclude category filtering resulted in no data in the data + frame. Accepts "all" if you want to fallback to just including all categories. + + patterns (list of regex str): + Use this option to specify how the function/file/line are extracted from labels. + For example, the default patterns are: + + r"(?P.*) \\[(?P\\S+):(?P[0-9]+)\\]$" + r"(?P.*) \\[(?P\\S+)\\]$" + r"^(?P\\S+):(?P[0-9]+)$" + + To extract the function/file/line from the labeling patterns: + + func [file:line] + func [file] + file:line + + respectively. + + thread_index_regex (regex str): + Use this option to determine thread IDs from the name of the threads. For example, + Omnitrace labels certain process tracks "Thread X (S)" to indicate that the given + process track are samples of Thread X. The default pattern is: + + r"(T|t)hread (?P[0-9]+)( |$)" + + to extract the "thread_index" field. + + max_depth (int): + Set this a positive non-zero number signifying the maximum call-stack depth to + process. This can significantly reduce the processing time for large traces. + """ + + self.filename = filename if isinstance(filename, (list, tuple)) else [filename] + self.metadata = {"hatchet_inclusive_suffix": ".inc"} + self.default_metric = "time{}".format(self.metadata["hatchet_inclusive_suffix"]) + self.verbose = 0 + self.report = [] + self.exclude = [] + self.include = [] + self.categories = [] + self.default_categories = [] + self.df_categories = [] + self.dataframe = pd.DataFrame() + self.process = pd.DataFrame() + self.threads = pd.DataFrame() + self.track_ids = [] + self.trace_processor = [] + self.compiled_patterns = [] + self.thread_index_regex = None + self.max_depth = None + self.configure(**kwargs) + + def configure(self, **kwargs): + + # pre-compile the regex patterns for extracting the func, file, and line info + # users can use their own pattens via patterns=[...]. An empty set of patterns + # is valid for avoiding parsing func/file/line info + _default_patterns = [ + # func [file:line] + r"(?P.*) \[(?P\S+):(?P[0-9]+)\]$", + # func [file] + r"(?P.*) \[(?P\S+)\]$", + # file:line + r"^(?P\S+):(?P[0-9]+)$", + ] + _patterns = kwargs["patterns"] if "patterns" in kwargs else None + + if _patterns is None: + _patterns = _default_patterns + elif kwargs.get("use_default_patterns", True): + _patterns += _default_patterns + + self.compiled_patterns = [re.compile(x) for x in _patterns] + self.thread_index_regex = re.compile( + kwargs.get("thread_index_regex", "(T|t)hread (?P[0-9]+)( |$)") + ) + + def report_at_verbosity(key, lvl): + if self.verbose >= lvl and key not in self.report: + self.report.append(key) + + self.verbose = self.verbose if "verbose" not in kwargs else kwargs["verbose"] + self.report = kwargs["report"] if "report" in kwargs else self.report + report_at_verbosity("category", 1) + report_at_verbosity("process", 2) + report_at_verbosity("threads", 2) + report_at_verbosity("track_ids", 2) + + _filenames = sorted(self.filename) + self.filename = kwargs.get("filename", sorted(self.filename)) + + _new_filenames = [x for x in self.filename if x not in _filenames] + + if len(self.filename) + len(_new_filenames) != len(self.trace_processor): + self.trace_processor = [TraceProcessor(trace=(f)) for f in self.filename] + elif _new_filenames: + self.trace_processor += [TraceProcessor(trace=(f)) for f in _new_filenames] + + self.max_depth = kwargs.get("max_depth", None) + + def query_tp(self, query, index_name=lambda x: "tp_index"): + """Simplifies querying the trace processor and always adds a + "tp_index" column for referencing which trace_processor + generated the data + """ + + def _append_column(df, name, idx): + """Used to add a tp_index column to our data which is used to identify + which trace-processor the query results came from""" + if name and name not in df: + df.insert(0, name, idx) + return df + + def _get_dataframe(tp): + """workaround for bug in TraceProcessor.QueryResultIterator.as_pandas_dataframe()""" + query_itr = tp.query(f"{query}") + # the perfetto trace processor query function looks like this: + # + # def query(self, sql: str): + # response = self.http.execute_query(sql) + # if response.error: + # raise TraceProcessorException(response.error) + # return TraceProcessor.QueryResultIterator(response.column_names, + # response.batch) + # + # unfortunately, data type of response.column_names is RepeatedScalarContainer + # and in a lot of versions of pandas, this type does not satisfy any of it's + # checks for whether this is a valid Index-type for the columns: + # isinstance(..., Index) + # isinstance(..., ABCSeries) + # is_iterator(...) + # isinstance(..., list) + # + # and thus as_pandas_dataframe() raises a TypeError exception. Queries can + # be VERY expensive for a large database (>> 10s of seconds) so instead of + # try -> except -> re-query (necessary) -> convert to list -> re-call + # as_pandas_dataframe() like so: + # + # try: + # return query_itr.as_pandas_dataframe() + # except TypeError: + # query_itr = tp.query(...) + # query_itr.__dict__["..."] = list(query_itr.__dict__["..."]) + # return query_itr.as_pandas_dataframe() + # + # which would effectively result in two queries every single time, we just + # do it upfront if the dict entry that is known to cause the problem exists + # and is of the type that we know causes problems + # + _buggy_dict_entry = "_QueryResultIterator__column_names" + if ( + _buggy_dict_entry in query_itr.__dict__ + and type(query_itr.__dict__[_buggy_dict_entry]).__name__ + == "RepeatedScalarContainer" + ): + query_itr.__dict__[_buggy_dict_entry] = list( + query_itr.__dict__[_buggy_dict_entry] + ) + + return query_itr.as_pandas_dataframe() + + return pd.concat( + [ + _append_column( + _get_dataframe(tp), + index_name(idx), + idx, + ) + for idx, tp in enumerate(self.trace_processor) + ] + ) + + def extract_tp_data(self, **kwargs): + """Extracts all the necessary data from the trace processor""" + self.configure(**kwargs) + + self.dataframe = self.query_tp( + "SELECT slice_id, track_id, category, depth, stack_id, parent_stack_id, ts, dur, name FROM slice" + ) + + self.df_categories = sorted(list(self.dataframe["category"].unique())) + + # check for update to include/exclude category + self.exclude = kwargs.get("exclude_category", self.exclude) + self.include = kwargs.get("include_category", self.include) + + self.categories = self.df_categories[:] + + # apply include first + if self.include: + self.categories = [x for x in self.categories if x in self.include] + + # apply exclude after + if self.exclude: + self.categories = [x for x in self.categories if x not in self.exclude] + + self.default_categories = kwargs.get( + "default_categories", self.default_categories + ) + _acceptable_default_categories = ( + 'default_categories can be set to: "all", ["all"], or [list of categories...]' + ) + + if not self.categories and self.default_categories: + if not isinstance(self.default_categories, (tuple, list)): + self.default_categories = [self.default_categories] + + if "all" in self.default_categories: + self.categories = self.df_categories[:] + else: + raise ValueError( + f"invalid default_categories value: {self.default_categories}. {_acceptable_default_categories}" + ) + + # filter out any categories that do not exist + self.categories = sorted([x for x in self.categories if x in self.df_categories]) + + if not self.categories: + raise ValueError( + f"The application of include_category={self.include} followed by exclude_category={self.exclude} rendered an empty set of categories (available={self.df_categories}). Either clear one of the configs or assign the default_categories. {_acceptable_default_categories}" + ) + + if "category" in self.report: + _ignore = [x for x in self.df_categories if x not in self.categories] + print( + "categories: {}{}".format( + ", ".join(self.categories), + " (ignored: {})".format(", ".join(_ignore)) if _ignore else "", + ) + ) + + # reduce the dataframe to given specified category data + # TODO: adjust the parent stack ids. if category entry is child of category entry, we lose category entry + self.dataframe = self.dataframe[self.dataframe["category"].isin(self.categories)] + + if self.dataframe.empty: + raise RuntimeError( + "category filtering resulted in an empty dataframe. categories: include={}, exclude={}, available={}".format( + self.include, self.exclude, self.df_categories + ) + ) + + self.process = self.query_tp( + "SELECT process.upid AS process_upid, process.id AS process_id, process.pid, process.name AS process_name, process_track.upid as track_upid, process_track.id AS track_id, process_track.parent_id as track_parent_id, process_track.name AS track_name from process JOIN process_track ON process_track.upid = process.upid WHERE process.pid > 0" + ) + self.threads = self.query_tp( + "SELECT thread.utid AS thread_utid, thread.id AS thread_id, thread.tid, thread.name as thread_name, thread.is_main_thread, thread_track.id AS track_id, thread_track.parent_id AS track_parent_id, thread_track.name AS track_name from thread JOIN thread_track ON thread_track.utid = thread.utid" + ) + + # generate empty dictionaries for each trace processor + self.track_ids = [{} for _ in range(len(self.trace_processor))] + + # generate mapping from track IDs to process and thread info. + # the "pid" and "tid" fields are the system value. we want to + # assign a "rank" and "thread" value for "pid" and "tid", + # respectively which start at zero and monotonically increase + for thread in self.threads.itertuples(): + _thread_name = ( + thread.thread_name if thread.track_name is None else thread.track_name + ) + for process in self.process.itertuples(): + if process.tp_index != thread.tp_index: + continue + _process_name = ( + process.process_name + if process.track_name is None + else process.track_name + ) + if process.track_id == thread.track_parent_id: + self.track_ids[thread.tp_index][thread.track_id] = { + "tp_index": thread.tp_index, + "pid": process.pid, + "tid": thread.tid, + "rank": -1, + "thread": -1, + "prio": 0 if thread.is_main_thread else 1, + "process_name": _process_name, + "thread_name": _thread_name, + } + break + + # some track ids do not have an associated system thread id so handle them here. + # for example, omnitrace post-processes sampling data collected on a thread + # during finalization and is inserted into perfetto on the main thread + # but not in the main thread track so perfetto does not associate the + # track_id with a system thread id. + for process in self.process.itertuples(): + if process.track_id in self.track_ids[process.tp_index].keys(): + continue + _process_name = ( + process.track_name + if process.process_name is None + else process.process_name + ) + _thread_name = ( + process.process_name if process.track_name is None else process.track_name + ) + self.track_ids[process.tp_index][process.track_id] = { + "tp_index": process.tp_index, + "pid": process.pid, + "tid": process.pid, + "rank": -1, + "thread": -1, + "prio": 0 if process.track_parent_id is None else 2, + "process_name": _process_name, + "thread_name": _thread_name, + } + + if "track_ids" in self.report and self.verbose >= 3: + print("\ntrack ids (original):") + for idx, _track_ids in enumerate(self.track_ids): + for key, itr in _track_ids.items(): + print(f" {idx:2}: {key:8} :: {itr}") + print("") + + # since the protobuf just has raw (system) PID and TIDs and there may be multiple PIDs and TIDs + # in the same file, we need to map the system PIDs to rank IDs starting at zero and, for each + # PID, map the system TIDs to thread-ids starting at zero + pid_offset = 0 + for idx, _track_ids in enumerate(self.track_ids): + _track_ids = dict( + sorted( + _track_ids.items(), + key=lambda x: [x[1]["pid"], x[1]["prio"], x[1]["tid"]], + ) + ) + pids = list(set([x["pid"] for _, x in _track_ids.items()])) + + if self.verbose >= 3: + tids = list(set([x["tid"] for _, x in _track_ids.items()])) + print(f"pids: {pids}") + print(f"tids: {tids}") + + # assign the rank and then increment the rank offset by the number of PIDs in the file + for pidx, pid in enumerate(pids): + for _, itr in _track_ids.items(): + if itr["pid"] == pid: + itr["rank"] = pidx + pid_offset + pid_offset += len(pids) + + for _, pid in enumerate(pids): + # dictionary containing only the data for this pid + _pid_track_ids = dict( + [[x, y] for x, y in _track_ids.items() if y["pid"] == pid] + ) + + assigned_track_ids = [] + + # filter out the main threads (priority == 0) for a given pid and set index to a value of zero + main_thr_info = set( + [x for x, y in _pid_track_ids.items() if y["prio"] == 0] + ) + + # for known "main" threads, assign index to zero + for track_id, track_id_data in _track_ids.items(): + if track_id in main_thr_info: + track_id_data["thread"] = 0 + assigned_track_ids.append(track_id) + + # starting value for assignment. set before next step + offset = 1 if assigned_track_ids else 0 + + # search thread name to try to identify which thread it belongs to. + # needs to come after offset assignment + for track_id, track_id_data in _track_ids.items(): + if ( + track_id in assigned_track_ids + or track_id not in _pid_track_ids.keys() + ): + continue + m = re.search(self.thread_index_regex, track_id_data["thread_name"]) + if m: + track_id_data["thread"] = int(m.group("thread_index")) + assigned_track_ids.append(track_id) + + # filter out the non-main threads (priority > 0) for a given pid that haven't already been assigned an index + chld_thr_info = set( + [ + x + for x, y in _pid_track_ids.items() + if y["prio"] > 0 and x not in assigned_track_ids + ] + ) + + # finally, assign remaining tracks thread indexes via incrementing offset value + for track_id, track_id_data in _track_ids.items(): + if ( + track_id in assigned_track_ids + or track_id not in _pid_track_ids.keys() + ): + continue + if track_id in chld_thr_info: + track_id_data["thread"] = offset + assigned_track_ids.append(track_id) + offset += 1 + + # make sure the thread indexes are monotonically increasing + # this may not be the case because of the assignment via regex matching + _pid_track_ids = dict( + [[x, y] for x, y in _track_ids.items() if y["pid"] == pid] + ) + tidx_max = max([y["thread"] for x, y in _pid_track_ids.items()]) + tidx_uniq = len( + set( + [ + y["thread"] + for x, y in _pid_track_ids.items() + if y["thread"] >= 0 + ] + ) + ) + + if self.verbose >= 3: + print(f"\nTID :: max={tidx_max}, unique={tidx_uniq}\n") + + # add one to comparison since one thread with a value of 0 would be a size of 1 + while tidx_max + 1 > tidx_uniq: + for idx in range(tidx_max): + # if this is empty, we need to decrement all thread indexes > idx + _tidx_loc = [ + x for x, y in _pid_track_ids.items() if y["thread"] == idx + ] + if not _tidx_loc: + for itr in _pid_track_ids.keys(): + if _track_ids[itr]["thread"] > idx: + _track_ids[itr]["thread"] -= 1 + break + # exit the loop so that we recalculate tidx_max + tidx_max = max([y["thread"] for x, y in _pid_track_ids.items()]) + + if "process" in self.report: + print(f"\nprocess:\n{self.process.to_string()}\n") + if "threads" in self.report: + print(f"\nthreads:\n{self.threads.to_string()}\n") + if "track_ids" in self.report: + print("\ntrack ids:") + for idx, _track_ids in enumerate(self.track_ids): + for key, itr in _track_ids.items(): + print(f" {idx:2}: {key:8} :: {itr}") + print("") + + if self.verbose >= 3: + print("\nTID mapping:") + for idx, _track_ids in enumerate(self.track_ids): + for track_id, itr in _track_ids.items(): + pid = itr["pid"] + tid = itr["tid"] + pidx = itr["rank"] + tidx = itr["thread"] + print( + f" {idx:2}: [{track_id:4}] {pid:8} -> {pidx:8} :: {tid:8} -> {tidx:8}" + ) + + print("") + + def create_graph(self): + """Create graph and dataframe""" + + def get_frame_attributes(_name): + """Get the standard set of dictionary entries for a Frame. + Also, parses the prefix for func-file-line info + which is typically in the form: + [:] + [] + """ + + if not self.compiled_patterns: + return {"type": "function", "name": _name} + + def _process_regex(_data): + """Process the regex data for func/file/line info""" + return _data.groupdict() if _data is not None else None + + def _perform_regex(_name): + """Performs a search for standard configurations of function + file + line""" + for _pattern in self.compiled_patterns: + _tmp = _process_regex(re.search(_pattern, _name)) + if _tmp: + return _tmp + return None + + _keys = {"type": "region", "name": _name} + _extra = {"file": "", "line": "0"} + _pdict = _perform_regex(_name) + if _pdict is not None: + _func = _pdict.get("func", None) + _file = _pdict.get("file", "") + _line = _pdict.get("line", "0") + _head = _pdict.get("head", None) + _tail = _pdict.get("tail", None) + + _line_s = f":{_line}" if int(_line) > 0 else "" + _tail_s = f"/{_tail}" if _tail is not None else "" + _file_s = f"{_file}{_line_s}" if _file != "" else _file + + _extra["file"] = _file_s + _extra["line"] = _line + + if "head" in _pdict: + _keys["name"] = _head.rstrip() + if _func is not None: + _extra["func"] = _func + else: + if _func is not None: + _keys["name"] = _func + else: + _keys["name"] = _file_s + + _keys["name"] = "{}{}".format(_keys["name"], _tail_s) + + return (_keys, _extra) + + # list_roots = [] + track_id_dict = OrderedDict() + callpath_to_node = {} + + def_metric = self.default_metric + + df = self.dataframe + _cols = [ + "tp_index", + "track_id", + "category", + "slice_id", + "stack_id", + "parent_stack_id", + "name", + "depth", + "ts", + "dur", + ] + _data = [df[x].to_list() for x in _cols] + + assert min([len(x) for x in _data]) == max([len(x) for x in _data]) + + for _tp_index, _track_id in zip( + _data[_cols.index("tp_index")], _data[_cols.index("track_id")] + ): + assert _tp_index < len(self.track_ids) + assert _track_id in self.track_ids[_tp_index].keys() + + _track_info = self.track_ids[_tp_index][_track_id] + _rank = _track_info["rank"] + _thread = _track_info["thread"] + + if _tp_index not in track_id_dict: + track_id_dict[_tp_index] = OrderedDict() + if _rank not in track_id_dict[_tp_index]: + track_id_dict[_tp_index][_rank] = OrderedDict() + if _thread not in track_id_dict[_tp_index][_rank]: + track_id_dict[_tp_index][_rank][_thread] = OrderedDict() + + track_id_dict[_tp_index][_rank][_thread] = {0: None} + + for ( + _tp_index, + _track_id, + _category, + _slice_id, + _stack_id, + _parent_stack_id, + _name, + _depth, + _ts, + _dur, + ) in zip(*_data): + _track_info = self.track_ids[_tp_index][_track_id] + _rank = _track_info["rank"] + _thread = _track_info["thread"] + + _track_id_dict = track_id_dict[_tp_index][_rank][_thread] + + # removed because of filtering + if _parent_stack_id not in _track_id_dict: + continue + + # reduce processing time + if self.max_depth is not None and _depth > self.max_depth: + continue + + _metrics = {} + _metrics["rank"] = _track_info["rank"] + _metrics["thread"] = _track_info["thread"] + _metrics["pid"] = _track_info["pid"] + _metrics["tid"] = _track_info["tid"] + _metrics["track_id"] = _track_id + _metrics["slice_id"] = _slice_id + _metrics["stack_id"] = _stack_id + _metrics["parent_stack_id"] = _parent_stack_id + _metrics["ts"] = _ts + _metrics[def_metric] = float(_dur) * 1.0e-9 # nsec -> sec + + _frame_attrs, _extra = get_frame_attributes(_name) + _extra["tp_index"] = _tp_index + _extra["category"] = _category + _extra["depth"] = _depth + + # look up the parent node specific to the TP index, rank, and thread + # stack ID is assigned by perfetto and parent stack ID is the + # stack ID of it's parent. + # _parent_node = _track_id_dict[_parent_stack_id] + + # hnode = Node(Frame(_frame_attrs, **_extra), None) + + # if _parent_node: + # _parent_node.add_child(hnode) + # else: + # list_roots.append(hnode) + + # # make sure this stack ID is unique for the + # # TP index, rank, and thread and is equal to a + # # previously seen node with the same stack ID + # if _stack_id not in _track_id_dict: + # _track_id_dict[_stack_id] = hnode + # elif _track_id_dict[_stack_id] != hnode: + # _existing = _track_id_dict[_stack_id] + # raise RuntimeError( + # f"{_stack_id} already exists in track_id_dict[{_tp_index}][{_rank}][{_thread}]. failed to set:\n {hnode.frame} (current)\n {_existing.frame} (existing)" + # ) + + _hash = hash((_tp_index, _slice_id)) + if _hash in callpath_to_node: + raise ValueError(f"{_hash} already exists in callpath_to_node dict") + + _frame_attrs.pop("type") # should not be a column in dataframe + callpath_to_node[_hash] = dict( + **_frame_attrs, + **_metrics, + **_extra, + ) + + if not callpath_to_node: + raise RuntimeError( + "call-graph is empty. if category filtering was used, you may have filtered out all the root nodes and thus all of it's children" + ) + + return (df, callpath_to_node) + # graph, dataframe = graphframe_indexing_helper( + # list_roots, + # data=list(callpath_to_node.values()), + # extensions=["rank", "thread"], + # fill_value=0, + # ) + + # inc_metrics = [self.default_metric] + # exc_metrics = [] + # return (graph, dataframe, exc_metrics, inc_metrics, self.default_metric) + + def read(self, **kwargs): + """Read perfetto json.""" + + self.configure(**kwargs) + self.extract_tp_data(**kwargs) + + ( + dataframe, + callpath_to_node, + ) = self.create_graph() + + # def _read(**kwargs): + # return self.read(**kwargs) + + # def _configure(**kwargs): + # self.configure(**kwargs) + + # def _query(*args, **kwargs): + # assert self.trace_processor + # return self.query_tp(*args, **kwargs) + + return ( + dataframe, + callpath_to_node, + ) + # default_metric=def_metric, + # metadata=self.metadata, + # attributes={ + # "reader": self, + # "read": _read, + # "query": _query, + # "configure": _configure, + # "selected_categories": lambda: self.categories, + # "available_categories": lambda: self.df_categories, + # }, + # ) diff --git a/tests/pytest-packages/tests/__init__.py b/tests/pytest-packages/tests/__init__.py new file mode 100644 index 0000000000..24d26e36c3 --- /dev/null +++ b/tests/pytest-packages/tests/__init__.py @@ -0,0 +1,23 @@ +# MIT License +# +# Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to deal +# in the Software without restriction, including without limitation the rights +# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +# copies of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in all +# copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +# SOFTWARE. + +from __future__ import absolute_import diff --git a/tests/pytest-packages/tests/rocprofv3.py b/tests/pytest-packages/tests/rocprofv3.py new file mode 100644 index 0000000000..53d776afb8 --- /dev/null +++ b/tests/pytest-packages/tests/rocprofv3.py @@ -0,0 +1,48 @@ +# MIT License +# +# Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to deal +# in the Software without restriction, including without limitation the rights +# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +# copies of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in all +# copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +# SOFTWARE. + +from __future__ import absolute_import + + +def test_perfetto_data( + pftrace_data, json_data, categories=("hip", "hsa", "marker", "kernel", "memory_copy") +): + + mapping = { + "hip": ("hip_api", "hip_api"), + "hsa": ("hsa_api", "hsa_api"), + "marker": ("marker_api", "marker_api"), + "kernel": ("kernel_dispatch", "kernel_dispatches"), + "memory_copy": ("memory_copy", "memory_copy"), + } + + # make sure they specified valid categories + for itr in categories: + assert itr in mapping.keys() + + for pf_category, js_category in [ + itr for key, itr in mapping.items() if key in categories + ]: + _pf_data = pftrace_data.loc[pftrace_data["category"] == pf_category] + _js_data = json_data["rocprofiler-sdk-tool"]["buffer_records"][js_category] + + assert len(_pf_data) == len(_js_data) diff --git a/tests/rocprofv3/hsa-queue-dependency/CMakeLists.txt b/tests/rocprofv3/hsa-queue-dependency/CMakeLists.txt index 1de09e4926..26ebcb1040 100644 --- a/tests/rocprofv3/hsa-queue-dependency/CMakeLists.txt +++ b/tests/rocprofv3/hsa-queue-dependency/CMakeLists.txt @@ -22,7 +22,7 @@ add_test( NAME rocprofv3-test-hsa-multiqueue-execute COMMAND $ --hsa-trace --kernel-trace -d - ${CMAKE_CURRENT_BINARY_DIR}/%argt%-trace -o out --output-format json:csv + ${CMAKE_CURRENT_BINARY_DIR}/%argt%-trace -o out --output-format json:csv:pftrace $) set_tests_properties( @@ -38,7 +38,9 @@ add_test( --kernel-trace-input ${CMAKE_CURRENT_BINARY_DIR}/multiqueue_testapp-trace/out_kernel_trace.csv --json-input - ${CMAKE_CURRENT_BINARY_DIR}/multiqueue_testapp-trace/out_results.json) + ${CMAKE_CURRENT_BINARY_DIR}/multiqueue_testapp-trace/out_results.json + --pftrace-input + ${CMAKE_CURRENT_BINARY_DIR}/multiqueue_testapp-trace/out_results.pftrace) set(MULTIQUEUE_VALIDATION_FILES ${CMAKE_CURRENT_BINARY_DIR}/multiqueue_testapp-trace/out_hsa_api_trace.csv diff --git a/tests/rocprofv3/hsa-queue-dependency/conftest.py b/tests/rocprofv3/hsa-queue-dependency/conftest.py index 456e224bc5..f9e31ebd77 100644 --- a/tests/rocprofv3/hsa-queue-dependency/conftest.py +++ b/tests/rocprofv3/hsa-queue-dependency/conftest.py @@ -6,6 +6,7 @@ import json from rocprofiler_sdk.pytest_utils.dotdict import dotdict from rocprofiler_sdk.pytest_utils import collapse_dict_list +from rocprofiler_sdk.pytest_utils.perfetto_reader import PerfettoReader def pytest_addoption(parser): @@ -24,6 +25,11 @@ def pytest_addoption(parser): action="store", help="Path to JSON file.", ) + parser.addoption( + "--pftrace-input", + action="store", + help="Path to Perfetto trace file.", + ) @pytest.fixture @@ -55,3 +61,9 @@ def json_data(request): filename = request.config.getoption("--json-input") with open(filename, "r") as inp: return dotdict(collapse_dict_list(json.load(inp))) + + +@pytest.fixture +def pftrace_data(request): + filename = request.config.getoption("--pftrace-input") + return PerfettoReader(filename).read()[0] diff --git a/tests/rocprofv3/hsa-queue-dependency/validate.py b/tests/rocprofv3/hsa-queue-dependency/validate.py index bb26c341c4..d2985309a5 100644 --- a/tests/rocprofv3/hsa-queue-dependency/validate.py +++ b/tests/rocprofv3/hsa-queue-dependency/validate.py @@ -154,6 +154,14 @@ def test_hsa_api_trace_json(json_data): assert functions.count("hsa_signal_destroy") == num_hsa_signal_destroy_calls +def test_perfetto_data(pftrace_data, json_data): + import rocprofiler_sdk.tests.rocprofv3 as rocprofv3 + + rocprofv3.test_perfetto_data( + pftrace_data, json_data, ("hip", "hsa", "marker", "kernel", "memory_copy") + ) + + if __name__ == "__main__": exit_code = pytest.main(["-x", __file__] + sys.argv[1:]) sys.exit(exit_code) diff --git a/tests/rocprofv3/tracing-hip-in-libraries/CMakeLists.txt b/tests/rocprofv3/tracing-hip-in-libraries/CMakeLists.txt index cc6771a025..c3149faead 100644 --- a/tests/rocprofv3/tracing-hip-in-libraries/CMakeLists.txt +++ b/tests/rocprofv3/tracing-hip-in-libraries/CMakeLists.txt @@ -27,6 +27,15 @@ add_test( --hsa-finalizer-trace --kernel-trace --memory-copy-trace --output-format JSON -d ${CMAKE_CURRENT_BINARY_DIR}/%argt%-trace -o out $) +add_test( + NAME rocprofv3-test-trace-hip-in-libraries-pftrace-execute + COMMAND + $ --hip-runtime-trace + --hip-compiler-trace --hsa-core-trace --hsa-amd-trace --hsa-image-trace + --hsa-finalizer-trace --kernel-trace --memory-copy-trace --output-format pftrace + -d ${CMAKE_CURRENT_BINARY_DIR}/%argt%-trace -o out + $) + string(REPLACE "LD_PRELOAD=" "ROCPROF_PRELOAD=" PRELOAD_ENV "${ROCPROFILER_MEMCHECK_PRELOAD_ENV}") @@ -58,6 +67,19 @@ set_tests_properties( "HSA_CORE_API|HSA_AMD_EXT_API|HSA_IMAGE_EXT_API|HSA_FINALIZER_EXT_API|HIP_API|HIP_COMPILER_API|KERNEL_DISPATCH|CODE_OBJECT" ) +set_tests_properties( + rocprofv3-test-trace-hip-in-libraries-pftrace-execute + PROPERTIES + TIMEOUT + 100 + LABELS + "integration-tests" + ENVIRONMENT + "${tracing-env}" + FAIL_REGULAR_EXPRESSION + "HSA_CORE_API|HSA_AMD_EXT_API|HSA_IMAGE_EXT_API|HSA_FINALIZER_EXT_API|HIP_API|HIP_COMPILER_API|KERNEL_DISPATCH|CODE_OBJECT" + ) + rocprofiler_configure_pytest_files(CONFIG pytest.ini COPY validate.py conftest.py) add_test( @@ -79,9 +101,13 @@ add_test( --hsa-stats ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_hsa_stats.csv --memory-copy-stats ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_memory_copy_stats.csv - --json-input ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_results.json) + --json-input ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_results.json + --pftrace-input + ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_results.pftrace) set(VALIDATION_FILES + ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_results.pftrace + ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_results.json ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_memory_copy_trace.csv ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_hsa_api_trace.csv ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_hip_api_trace.csv @@ -90,8 +116,7 @@ set(VALIDATION_FILES ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_kernel_stats.csv ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_hip_stats.csv ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_hsa_stats.csv - ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_memory_copy_stats.csv - ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_results.json) + ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_memory_copy_stats.csv) set_tests_properties( rocprofv3-test-trace-hip-in-libraries-validate @@ -101,7 +126,7 @@ set_tests_properties( LABELS "integration-tests" DEPENDS - "rocprofv3-test-trace-hip-in-libraries-execute;rocprofv3-test-trace-hip-in-libraries-json-execute" + "rocprofv3-test-trace-hip-in-libraries-execute;rocprofv3-test-trace-hip-in-libraries-json-execute;rocprofv3-test-trace-hip-in-libraries-pftrace-execute" FAIL_REGULAR_EXPRESSION "AssertionError" ATTACHED_FILES_ON_FAIL diff --git a/tests/rocprofv3/tracing-hip-in-libraries/conftest.py b/tests/rocprofv3/tracing-hip-in-libraries/conftest.py index d6835b7488..2d1bef6ab8 100644 --- a/tests/rocprofv3/tracing-hip-in-libraries/conftest.py +++ b/tests/rocprofv3/tracing-hip-in-libraries/conftest.py @@ -7,6 +7,7 @@ import json from rocprofiler_sdk.pytest_utils.dotdict import dotdict from rocprofiler_sdk.pytest_utils import collapse_dict_list +from rocprofiler_sdk.pytest_utils.perfetto_reader import PerfettoReader def pytest_addoption(parser): @@ -66,6 +67,11 @@ def pytest_addoption(parser): action="store", help="Path to JSON file.", ) + parser.addoption( + "--pftrace-input", + action="store", + help="Path to Perfetto trace file.", + ) @pytest.fixture @@ -198,3 +204,9 @@ def json_data(request): filename = request.config.getoption("--json-input") with open(filename, "r") as inp: return dotdict(collapse_dict_list(json.load(inp))) + + +@pytest.fixture +def pftrace_data(request): + filename = request.config.getoption("--pftrace-input") + return PerfettoReader(filename).read()[0] diff --git a/tests/rocprofv3/tracing-hip-in-libraries/validate.py b/tests/rocprofv3/tracing-hip-in-libraries/validate.py index 775442ed36..cfe8b8b6b3 100644 --- a/tests/rocprofv3/tracing-hip-in-libraries/validate.py +++ b/tests/rocprofv3/tracing-hip-in-libraries/validate.py @@ -398,6 +398,15 @@ def test_memory_copy_trace( validate_stats(row) +def test_perfetto_data(pftrace_data, json_data): + import rocprofiler_sdk.tests.rocprofv3 as rocprofv3 + + # do not test for HSA since that may vary b/t two separate runs + rocprofv3.test_perfetto_data( + pftrace_data, json_data, ("hip", "marker", "kernel", "memory_copy") + ) + + if __name__ == "__main__": exit_code = pytest.main(["-x", __file__] + sys.argv[1:]) sys.exit(exit_code) diff --git a/tests/rocprofv3/tracing-plus-cc/CMakeLists.txt b/tests/rocprofv3/tracing-plus-cc/CMakeLists.txt index 52cb110d4b..2bffbd3966 100644 --- a/tests/rocprofv3/tracing-plus-cc/CMakeLists.txt +++ b/tests/rocprofv3/tracing-plus-cc/CMakeLists.txt @@ -20,7 +20,7 @@ add_test( COMMAND $ --hsa-trace -i ${CMAKE_CURRENT_BINARY_DIR}/input.txt -d ${CMAKE_CURRENT_BINARY_DIR}/out_cc_trace - -o pmc3 --output-format JSON:CSV $) + -o pmc3 --output-format JSON:PFTRACE:CSV $) string(REPLACE "LD_PRELOAD=" "ROCPROF_PRELOAD=" PRELOAD_ENV "${ROCPROFILER_MEMCHECK_PRELOAD_ENV}") @@ -39,6 +39,8 @@ foreach(_DIR "pmc_1" "pmc_2" "pmc_3" "pmc_4") COMMAND ${Python3_EXECUTABLE} ${CMAKE_CURRENT_BINARY_DIR}/validate.py --json-input "${CMAKE_CURRENT_BINARY_DIR}/out_cc_trace/${_DIR}/pmc3_results.json" + --pftrace-input + "${CMAKE_CURRENT_BINARY_DIR}/out_cc_trace/${_DIR}/pmc3_results.pftrace" --hsa-input "${CMAKE_CURRENT_BINARY_DIR}/out_cc_trace/${_DIR}/pmc3_hsa_api_trace.csv" --agent-input diff --git a/tests/rocprofv3/tracing-plus-cc/conftest.py b/tests/rocprofv3/tracing-plus-cc/conftest.py index 309203403b..6225384d33 100644 --- a/tests/rocprofv3/tracing-plus-cc/conftest.py +++ b/tests/rocprofv3/tracing-plus-cc/conftest.py @@ -8,6 +8,7 @@ import pandas as pd from rocprofiler_sdk.pytest_utils.dotdict import dotdict from rocprofiler_sdk.pytest_utils import collapse_dict_list +from rocprofiler_sdk.pytest_utils.perfetto_reader import PerfettoReader def pytest_addoption(parser): @@ -46,6 +47,11 @@ def pytest_addoption(parser): action="store", help="Path to JSON file.", ) + parser.addoption( + "--pftrace-input", + action="store", + help="Path to Perfetto trace file.", + ) @pytest.fixture @@ -184,3 +190,9 @@ def json_data(request): filename = request.config.getoption("--json-input") with open(filename, "r") as inp: return dotdict(collapse_dict_list(json.load(inp))) + + +@pytest.fixture +def pftrace_data(request): + filename = request.config.getoption("--pftrace-input") + return PerfettoReader(filename).read()[0] diff --git a/tests/rocprofv3/tracing-plus-cc/validate.py b/tests/rocprofv3/tracing-plus-cc/validate.py index 5b6d2f2edb..dc4ec6dc5b 100644 --- a/tests/rocprofv3/tracing-plus-cc/validate.py +++ b/tests/rocprofv3/tracing-plus-cc/validate.py @@ -20,6 +20,14 @@ def test_validate_counter_collection_plus_tracing( assert len(counter_collection_data) > 0 +def test_perfetto_data(pftrace_data, json_data): + import rocprofiler_sdk.tests.rocprofv3 as rocprofv3 + + rocprofv3.test_perfetto_data( + pftrace_data, json_data, ("hip", "hsa", "marker", "kernel", "memory_copy") + ) + + if __name__ == "__main__": exit_code = pytest.main(["-x", __file__] + sys.argv[1:]) sys.exit(exit_code) diff --git a/tests/rocprofv3/tracing/CMakeLists.txt b/tests/rocprofv3/tracing/CMakeLists.txt index 4df924cd7f..eb94169ef5 100644 --- a/tests/rocprofv3/tracing/CMakeLists.txt +++ b/tests/rocprofv3/tracing/CMakeLists.txt @@ -15,7 +15,7 @@ add_test( COMMAND $ -M --hsa-trace --kernel-trace --memory-copy-trace --marker-trace -d ${CMAKE_CURRENT_BINARY_DIR}/%argt%-trace -o - out --output-format csv,json $) + out --output-format pftrace,csv,json $) string(REPLACE "LD_PRELOAD=" "ROCPROF_PRELOAD=" PRELOAD_ENV "${ROCPROFILER_MEMCHECK_PRELOAD_ENV}") @@ -57,15 +57,18 @@ add_test( ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_marker_api_trace.csv --agent-input ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_agent_info.csv - --json-input ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_results.json) + --json-input ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_results.json + --pftrace-input + ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_results.pftrace) set(VALIDATION_FILES + ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_results.pftrace + ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_results.json ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_memory_copy_trace.csv ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_hsa_api_trace.csv ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_kernel_trace.csv ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_marker_api_trace.csv - ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_agent_info.csv - ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_results.json) + ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-trace/out_agent_info.csv) set_tests_properties( rocprofv3-test-trace-validate @@ -87,8 +90,8 @@ add_test( NAME rocprofv3-test-systrace-execute COMMAND $ --sys-trace -d - ${CMAKE_CURRENT_BINARY_DIR}/%argt%-systrace -o out --output-format csv,json - $) + ${CMAKE_CURRENT_BINARY_DIR}/%argt%-systrace -o out --output-format + pftrace,csv,json $) set_tests_properties( rocprofv3-test-systrace-execute @@ -118,15 +121,18 @@ add_test( --agent-input ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-systrace/out_agent_info.csv --json-input - ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-systrace/out_results.json) + ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-systrace/out_results.json + --pftrace-input + ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-systrace/out_results.pftrace) set(SYS_VALIDATION_FILES + ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-systrace/out_results.pftrace + ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-systrace/out_results.json ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-systrace/out_memory_copy_trace.csv ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-systrace/out_hsa_api_trace.csv ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-systrace/out_kernel_trace.csv ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-systrace/out_marker_api_trace.csv - ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-systrace/out_agent_info.csv - ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-systrace/out_results.json) + ${CMAKE_CURRENT_BINARY_DIR}/simple-transpose-systrace/out_agent_info.csv) set_tests_properties( rocprofv3-test-systrace-validate diff --git a/tests/rocprofv3/tracing/conftest.py b/tests/rocprofv3/tracing/conftest.py index 3db9c1af3e..d5b632e464 100644 --- a/tests/rocprofv3/tracing/conftest.py +++ b/tests/rocprofv3/tracing/conftest.py @@ -6,6 +6,7 @@ import json from rocprofiler_sdk.pytest_utils.dotdict import dotdict from rocprofiler_sdk.pytest_utils import collapse_dict_list +from rocprofiler_sdk.pytest_utils.perfetto_reader import PerfettoReader def pytest_addoption(parser): @@ -44,6 +45,11 @@ def pytest_addoption(parser): action="store", help="Path to JSON file.", ) + parser.addoption( + "--pftrace-input", + action="store", + help="Path to Perfetto trace file.", + ) @pytest.fixture @@ -123,3 +129,9 @@ def json_data(request): filename = request.config.getoption("--json-input") with open(filename, "r") as inp: return dotdict(collapse_dict_list(json.load(inp))) + + +@pytest.fixture +def pftrace_data(request): + filename = request.config.getoption("--pftrace-input") + return PerfettoReader(filename).read()[0] diff --git a/tests/rocprofv3/tracing/validate.py b/tests/rocprofv3/tracing/validate.py index f21afe777e..11e116e121 100644 --- a/tests/rocprofv3/tracing/validate.py +++ b/tests/rocprofv3/tracing/validate.py @@ -263,6 +263,14 @@ def test_marker_api_trace_json(json_data): assert marker["end_timestamp"] >= marker["start_timestamp"] +def test_perfetto_data(pftrace_data, json_data): + import rocprofiler_sdk.tests.rocprofv3 as rocprofv3 + + rocprofv3.test_perfetto_data( + pftrace_data, json_data, ("hip", "hsa", "marker", "kernel", "memory_copy") + ) + + if __name__ == "__main__": exit_code = pytest.main(["-x", __file__] + sys.argv[1:]) sys.exit(exit_code) diff --git a/tests/tools/json-tool.cpp b/tests/tools/json-tool.cpp index fa99bef12b..83dd2246ff 100644 --- a/tests/tools/json-tool.cpp +++ b/tests/tools/json-tool.cpp @@ -1436,14 +1436,15 @@ write_json(call_stack_t* _call_stack) } { + namespace sdk = ::rocprofiler::sdk; using JSONOutputArchive = cereal::MinimalJSONOutputArchive; constexpr auto json_prec = 32; constexpr auto json_indent = JSONOutputArchive::Options::IndentChar::space; auto json_opts = JSONOutputArchive::Options{json_prec, json_indent, 1}; auto json_ar = JSONOutputArchive{*ofs, json_opts}; - auto buffer_names = rocprofiler::sdk::get_buffer_tracing_names(); - auto callbk_names = rocprofiler::sdk::get_callback_tracing_names(); + auto buffer_names = sdk::get_buffer_tracing_names(); + auto callbk_names = sdk::get_callback_tracing_names(); auto validate_page_migration = (page_migration_status != ROCPROFILER_STATUS_ERROR_INCOMPATIBLE_KERNEL); @@ -1655,8 +1656,10 @@ write_perfetto() } { - auto buffer_names = rocprofiler::sdk::get_buffer_tracing_names(); - auto callbk_name_info = rocprofiler::sdk::get_callback_tracing_names(); + namespace sdk = ::rocprofiler::sdk; + + auto buffer_names = sdk::get_buffer_tracing_names(); + auto callbk_name_info = sdk::get_callback_tracing_names(); for(auto itr : hsa_api_bf_records) { @@ -1671,7 +1674,7 @@ write_perfetto() }); if(ritr != hsa_api_cb_records.end()) _args = ritr->args; - TRACE_EVENT_BEGIN(rocprofiler::trait::name::value, + TRACE_EVENT_BEGIN(sdk::perfetto_category::name, ::perfetto::StaticString(name.data()), track, itr.start_timestamp, @@ -1688,9 +1691,9 @@ write_perfetto() itr.correlation_id.internal, [&](::perfetto::EventContext ctx) { for(const auto& aitr : _args) - add_perfetto_annotation(ctx, aitr.first, aitr.second); + sdk::add_perfetto_annotation(ctx, aitr.first, aitr.second); }); - TRACE_EVENT_END(rocprofiler::trait::name::value, + TRACE_EVENT_END(sdk::perfetto_category::name, track, itr.end_timestamp, "end_ns", @@ -1710,7 +1713,7 @@ write_perfetto() }); if(ritr != hip_api_cb_records.end()) _args = ritr->args; - TRACE_EVENT_BEGIN(rocprofiler::trait::name::value, + TRACE_EVENT_BEGIN(sdk::perfetto_category::name, ::perfetto::StaticString(name.data()), track, itr.start_timestamp, @@ -1727,9 +1730,9 @@ write_perfetto() itr.correlation_id.internal, [&](::perfetto::EventContext ctx) { for(const auto& aitr : _args) - add_perfetto_annotation(ctx, aitr.first, aitr.second); + sdk::add_perfetto_annotation(ctx, aitr.first, aitr.second); }); - TRACE_EVENT_END(rocprofiler::trait::name::value, + TRACE_EVENT_END(sdk::perfetto_category::name, track, itr.end_timestamp, "end_ns", @@ -1741,7 +1744,7 @@ write_perfetto() auto name = buffer_names.at(itr.kind, itr.operation); auto& track = agent_tracks.at(itr.dst_agent_id.handle); - TRACE_EVENT_BEGIN(rocprofiler::trait::name::value, + TRACE_EVENT_BEGIN(sdk::perfetto_category::name, ::perfetto::StaticString(name.data()), track, itr.start_timestamp, @@ -1758,7 +1761,7 @@ write_perfetto() agents_map.at(itr.dst_agent_id).logical_node_id, "copy_bytes", itr.bytes); - TRACE_EVENT_END(rocprofiler::trait::name::value, + TRACE_EVENT_END(sdk::perfetto_category::name, track, itr.end_timestamp, "end_ns", @@ -1787,34 +1790,33 @@ write_perfetto() demangled.emplace(name, demangle(name)); } - TRACE_EVENT_BEGIN( - rocprofiler::trait::name::value, - ::perfetto::StaticString(demangled.at(name).c_str()), - track, - itr.start_timestamp, - ::perfetto::Flow::ProcessScoped(itr.correlation_id.internal), - "begin_ns", - itr.start_timestamp, - "kind", - itr.kind, - "agent", - agents_map.at(info.agent_id).logical_node_id, - "corr_id", - itr.correlation_id.internal, - "queue", - info.queue_id.handle, - "kernel_id", - info.kernel_id, - "private_segment_size", - info.private_segment_size, - "group_segment_size", - info.group_segment_size, - "workgroup_size", - info.workgroup_size.x * info.workgroup_size.y * info.workgroup_size.z, - "grid_size", - info.grid_size.x * info.grid_size.y * info.grid_size.z); + TRACE_EVENT_BEGIN(sdk::perfetto_category::name, + ::perfetto::StaticString(demangled.at(name).c_str()), + track, + itr.start_timestamp, + ::perfetto::Flow::ProcessScoped(itr.correlation_id.internal), + "begin_ns", + itr.start_timestamp, + "kind", + itr.kind, + "agent", + agents_map.at(info.agent_id).logical_node_id, + "corr_id", + itr.correlation_id.internal, + "queue", + info.queue_id.handle, + "kernel_id", + info.kernel_id, + "private_segment_size", + info.private_segment_size, + "group_segment_size", + info.group_segment_size, + "workgroup_size", + info.workgroup_size.x * info.workgroup_size.y * info.workgroup_size.z, + "grid_size", + info.grid_size.x * info.grid_size.y * info.grid_size.z); - TRACE_EVENT_END(rocprofiler::trait::name::value, + TRACE_EVENT_END(sdk::perfetto_category::name, track, itr.end_timestamp, "end_ns",