Adding Perfetto support (#867)

* Perfetto submodule

* include/rocprofiler-sdk/cxx/perfetto.hpp

- adapted from tests/common/perfetto.hpp
- updated json-tool to use <rocprofiler-sdk/cxx/perfetto.hpp>

* 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
This commit is contained in:
Jonathan R. Madsen
2024-05-22 15:51:12 -05:00
gecommit door GitHub
bovenliggende 92b7326910
commit 957bb7a4e5
39 gewijzigde bestanden met toevoegingen van 2530 en 291 verwijderingen
+6 -1
Bestand weergeven
@@ -55,9 +55,14 @@ usage() {
echo -e "\t#${GREY} usage (with custom dir): rocprofv3 --hsa-trace -d <out_dir> -o <file_name> <executable>${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 <out_dir> <executable>${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 <executable>"
echo -e "\t#${GREY} (JSON + CSV): rocprofv3 --sys-trace --output-format JSON,CSV <executable>"
echo -e "\t#${GREY} (JSON + PFTRACE): rocprofv3 --sys-trace --output-format JSON,PFTRACE <executable>"
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 ""
@@ -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}
@@ -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}
@@ -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 <rocprofiler-sdk/cxx/details/mpl.hpp>
#include <functional>
#include <sstream>
#include <stdexcept>
#include <string>
#include <vector>
namespace rocprofiler
{
namespace sdk
{
namespace parse
{
template <typename Tp>
inline Tp
from_string(const std::string& str)
{
auto ss = std::stringstream{str};
auto val = Tp{};
ss >> val;
return val;
}
template <typename Tp>
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 ContainerT = std::vector<std::string>,
typename ValueT = typename ContainerT::value_type,
typename PredicateT = std::function<ValueT(ValueT&&)>>
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<PredicateT>(predicate)(std::move(_tmp)));
}
}
return _result;
}
/// \brief apply a string transformation to substring in between a common delimiter.
///
template <typename PredicateT = std::function<std::string(const std::string&)>>
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
@@ -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 <rocprofiler-sdk/cxx/details/mpl.hpp>
#include <array>
#include <cstring>
#include <initializer_list>
#include <ios>
#include <sstream>
#include <string>
#include <string_view>
#include <tuple>
#include <type_traits>
namespace rocprofiler
{
namespace sdk
{
namespace join
{
template <typename... ArgsT>
inline void
consume_args(ArgsT&&...)
{}
enum
{
NoQuoteStrings = 0x0,
QuoteStrings = 0x1
};
template <size_t Idx>
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 <int TraitT, typename ArgT>
inline decltype(auto)
join_arg(config _cfg, ArgT&& _v)
{
using arg_type = mpl::basic_identity_t<ArgT>;
constexpr bool _is_string_type = mpl::is_string_type<arg_type>::value;
constexpr bool _is_iterable = mpl::is_iterable<arg_type>(0);
constexpr bool _has_traits_type = mpl::has_traits<arg_type>(0);
constexpr bool _has_key_type = mpl::has_key_type<arg_type>(0);
constexpr bool _has_value_type = mpl::has_value_type<arg_type>(0);
constexpr bool _has_mapped_type = mpl::has_mapped_type<arg_type>(0);
if constexpr(_is_string_type)
{
if constexpr(TraitT == QuoteStrings)
{
return std::string{"\""} + std::string{std::forward<ArgT>(_v)} + std::string{"\""};
}
else
{
return std::forward<ArgT>(_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<ArgT>(_v))
_ss << _cfg.array.delimiter << _cfg.pair.prefix << join_arg<TraitT>(_cfg, itr.first)
<< _cfg.pair.delimiter << join_arg<TraitT>(_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<ArgT>(_v))
_ss << _cfg.array.delimiter << join_arg<TraitT>(_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<ArgT>(0))
{
return std::forward<ArgT>(_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<ArgT>::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 <int TraitT = NoQuoteStrings, typename... Args>
auto
join(config _cfg, Args&&... _args)
{
static_assert(std::is_trivially_copyable<config>::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<TraitT>(_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 <int TraitT = NoQuoteStrings, typename... Args>
auto
join(std::array<std::string_view, 3>&& _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>(_args)...);
}
template <int TraitT = NoQuoteStrings,
typename DelimT,
typename... Args,
std::enable_if_t<!mpl::is_basic_same<config, DelimT>::value, int> = 0>
auto
join(DelimT&& _delim, Args&&... _args)
{
using delim_type = mpl::basic_identity_t<DelimT>;
if constexpr(std::is_constructible<config, delim_type>::value)
{
auto _cfg = config{std::forward<DelimT>(_delim)};
return join<TraitT>(_cfg, std::forward<Args>(_args)...);
}
else if constexpr(std::is_same<delim_type, char>::value)
{
auto _cfg = config{};
const char _delim_c[2] = {_delim, '\0'};
_cfg.delimiter = _delim_c;
return join<TraitT>(_cfg, std::forward<Args>(_args)...);
}
else
{
auto _cfg = config{};
_cfg.delimiter = std::string_view{_delim};
return join<TraitT>(_cfg, std::forward<Args>(_args)...);
}
}
template <typename ArgT>
auto
quoted(ArgT&& _arg)
{
auto _cfg = config{};
_cfg.prefix = "\"";
_cfg.suffix = "\"";
return join(_cfg, std::forward<ArgT>(_arg));
}
} // namespace join
} // namespace sdk
} // namespace rocprofiler
@@ -23,17 +23,61 @@
#pragma once
#include <cstddef>
#include <cstdint>
#include <optional>
#include <string>
#include <string_view>
#include <type_traits>
#include <utility>
#define ROCPROFILER_IMPL_HAS_CONCEPT(NAME, TRAIT) \
template <typename Tp, typename = typename Tp::TRAIT> \
inline constexpr bool NAME(int) \
{ \
return true; \
} \
\
template <typename Tp> \
inline constexpr bool NAME(long) \
{ \
return false; \
}
#define ROCPROFILER_IMPL_SFINAE_CONCEPT(NAME, ...) \
template <typename Tp> \
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 <typename Tp>
struct unqualified_identity
{
using type = std::remove_cv_t<std::remove_reference_t<std::decay_t<Tp>>>;
};
template <typename Tp>
using unqualified_identity_t = typename unqualified_identity<Tp>::type;
template <typename Tp, typename Up>
struct is_same_unqualified_identity
: std::is_same<unqualified_identity_t<Tp>, unqualified_identity_t<Up>>
{};
template <typename Tp>
struct string_support
{
@@ -79,6 +123,135 @@ struct string_support<std::string>
type operator()(const char* val) const { return type{val}; }
};
namespace impl
{
template <typename Tp>
struct is_string_type : std::false_type
{};
template <>
struct is_string_type<std::string> : std::true_type
{};
template <>
struct is_string_type<char*> : std::true_type
{};
template <>
struct is_string_type<const char*> : std::true_type
{};
template <>
struct is_string_type<std::string_view> : std::true_type
{};
} // namespace impl
template <typename Tp>
struct is_string_type : impl::is_string_type<unqualified_identity_t<Tp>>
{};
// template <typename Tp>
// struct can_stringify
// {
// private:
// static constexpr auto sfinae(int)
// -> decltype(std::declval<std::ostream&>() << std::declval<Tp>(), 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<Tp>().empty())
ROCPROFILER_IMPL_SFINAE_CONCEPT(can_stringify, std::declval<std::ostream&>() << std::declval<Tp>())
ROCPROFILER_IMPL_SFINAE_CONCEPT(is_iterable,
std::begin(std::declval<Tp>()),
std::end(std::declval<Tp>()))
// compatability
template <typename Tp>
using supports_ostream = can_stringify<Tp>;
template <typename ArgT>
inline bool
is_empty(ArgT&& _v)
{
using arg_type = unqualified_identity_t<ArgT>;
if constexpr(has_empty_member_function<arg_type>::value)
{
return std::forward<ArgT>(_v).empty();
}
else if constexpr(is_string_type<arg_type>::value)
{
static_assert(std::is_constructible<std::string_view, ArgT>::value,
"not string_view constructible");
return std::string_view{std::forward<ArgT>(_v)}.empty();
}
return false;
}
namespace impl
{
template <typename ContainerT, typename... Args>
inline auto
emplace(ContainerT& _c, int, Args&&... _args)
-> decltype(_c.emplace_back(std::forward<Args>(_args)...))
{
return _c.emplace_back(std::forward<Args>(_args)...);
}
template <typename ContainerT, typename... Args>
inline auto
emplace(ContainerT& _c, long, Args&&... _args) -> decltype(_c.emplace(std::forward<Args>(_args)...))
{
return _c.emplace(std::forward<Args>(_args)...);
}
template <typename ContainerT, typename ArgT>
inline auto
reserve(ContainerT& _c, int, ArgT _arg) -> decltype(_c.reserve(_arg), bool())
{
_c.reserve(_arg);
return true;
}
template <typename ContainerT, typename ArgT>
inline auto
reserve(ContainerT&, long, ArgT)
{
return false;
}
} // namespace impl
template <typename ContainerT, typename... Args>
inline auto
emplace(ContainerT& _c, Args&&... _args)
{
return impl::emplace(_c, 0, std::forward<Args>(_args)...);
}
template <typename ContainerT, typename ArgT>
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
@@ -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 <typename Tp>
bool
equal(Tp lhs, Tp rhs) ROCPROFILER_ATTRIBUTE(pure);
template <typename Tp>
bool
less(Tp lhs, Tp rhs) ROCPROFILER_ATTRIBUTE(pure);
template <typename Tp>
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 <typename Tp>
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
@@ -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 <rocprofiler-sdk/cxx/details/mpl.hpp>
#include <cstddef>
#include <ostream>
#include <sstream>
#include <string>
#include <string_view>
#include <type_traits>
#include <utility>
#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 <typename Tp>
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.h>
PERFETTO_DEFINE_CATEGORIES(ROCPROFILER_PERFETTO_CATEGORIES);
namespace rocprofiler
{
namespace sdk
{
using perfetto_event_context_t = ::perfetto::EventContext;
template <typename Np, typename Tp>
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<Np>;
using value_type = mpl::unqualified_identity_t<Tp>;
static_assert(mpl::is_string_type<named_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<Np>(_name)}.data());
return _dbg;
};
if constexpr(std::is_same<value_type, std::string_view>::value)
{
_get_dbg()->set_string_value(_val.data());
}
else if constexpr(mpl::is_string_type<value_type>::value)
{
_get_dbg()->set_string_value(std::forward<Tp>(_val));
}
else if constexpr(std::is_same<value_type, bool>::value)
{
_get_dbg()->set_bool_value(_val);
}
else if constexpr(std::is_enum<value_type>::value)
{
_get_dbg()->set_int_value(static_cast<int64_t>(_val));
}
else if constexpr(std::is_floating_point<value_type>::value)
{
_get_dbg()->set_double_value(static_cast<double>(_val));
}
else if constexpr(std::is_integral<value_type>::value)
{
if constexpr(std::is_unsigned<value_type>::value)
{
_get_dbg()->set_uint_value(_val);
}
else
{
_get_dbg()->set_int_value(_val);
}
}
else if constexpr(std::is_pointer<value_type>::value)
{
_get_dbg()->set_pointer_value(reinterpret_cast<uint64_t>(_val));
}
else if constexpr(mpl::can_stringify<value_type>::value)
{
auto _ss = std::stringstream{};
_ss << std::forward<Tp>(_val);
_get_dbg()->set_string_value(_ss.str());
}
else
{
static_assert(std::is_empty<value_type>::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
@@ -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
@@ -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<std::string_view>{"CSV", "JSON"};
const auto supported_formats = std::set<std::string_view>{"CSV", "JSON", "PFTRACE"};
for(const auto& itr : entries)
{
LOG_IF(FATAL, supported_formats.count(itr) == 0)
@@ -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());
@@ -28,6 +28,7 @@
#include <rocprofiler-sdk/fwd.h>
#include <rocprofiler-sdk/marker/api_id.h>
#include <unistd.h>
#include <iomanip>
#include <string_view>
@@ -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 <rocprofiler-sdk/fwd.h>
#include <rocprofiler-sdk/marker/api_id.h>
#include <atomic>
#include <future>
#include <rocprofiler-sdk/cxx/hash.hpp>
#include <rocprofiler-sdk/cxx/operators.hpp>
#include <rocprofiler-sdk/cxx/perfetto.hpp>
#include <map>
#include <thread>
#include <unordered_map>
#include <utility>
namespace rocprofiler
{
namespace tool
{
namespace
{
auto main_tid = common::get_tid();
template <typename Tp>
size_t
get_hash_id(Tp&& _val)
{
if constexpr(!std::is_pointer<Tp>::value)
return std::hash<Tp>{}(std::forward<Tp>(_val));
else if constexpr(std::is_same<Tp, const char*>::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<rocprofiler_agent_v0_t> agent_data,
std::deque<rocprofiler_buffer_tracing_hip_api_record_t>* hip_api_data,
std::deque<rocprofiler_buffer_tracing_hsa_api_record_t>* hsa_api_data,
std::deque<rocprofiler_buffer_tracing_kernel_dispatch_record_t>* kernel_dispatch_data,
std::deque<rocprofiler_buffer_tracing_memory_copy_record_t>* memory_copy_data,
std::deque<rocprofiler_buffer_tracing_marker_api_record_t>* marker_api_data,
std::deque<rocprofiler_buffer_tracing_scratch_memory_record_t>* /*scratch_api_data*/)
{
namespace sdk = ::rocprofiler::sdk;
auto agents_map = std::unordered_map<rocprofiler_agent_id_t, rocprofiler_agent_t>{};
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<rocprofiler_thread_id_t>{};
auto demangled = std::unordered_map<std::string_view, std::string>{};
auto agent_thread_ids = std::unordered_map<rocprofiler_agent_id_t, std::set<uint64_t>>{};
auto agent_queue_ids =
std::unordered_map<rocprofiler_agent_id_t, std::unordered_set<rocprofiler_queue_id_t>>{};
auto thread_indexes = std::unordered_map<rocprofiler_thread_id_t, uint64_t>{};
auto kernel_sym_data = get_kernel_symbol_data();
auto thread_tracks = std::unordered_map<rocprofiler_thread_id_t, ::perfetto::Track>{};
auto agent_thread_tracks =
std::unordered_map<rocprofiler_agent_id_t,
std::unordered_map<uint64_t, ::perfetto::Track>>{};
auto agent_queue_tracks =
std::unordered_map<rocprofiler_agent_id_t,
std::unordered_map<rocprofiler_queue_id_t, ::perfetto::Track>>{};
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<sdk::category::hsa_api>::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<sdk::category::hsa_api>::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<sdk::category::hip_api>::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<sdk::category::hip_api>::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<sdk::category::marker_api>::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<sdk::category::marker_api>::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<sdk::category::memory_copy>::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<sdk::category::memory_copy>::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<sdk::category::kernel_dispatch>::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<sdk::category::kernel_dispatch>::name,
track,
itr.end_timestamp);
}
}
// counter tracks
{
// memory copy counter track
auto mem_cpy_endpoints = std::map<rocprofiler_agent_id_t, std::map<uint64_t, uint64_t>>{};
auto mem_cpy_extremes = std::pair<uint64_t, uint64_t>{};
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<rocprofiler_agent_id_t, ::perfetto::CounterTrack>{};
auto mem_cpy_cnt_names = std::vector<std::string>{};
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<sdk::category::memory_copy>::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<void(std::ostream*&)>{};
std::ostream* ofs = nullptr;
std::tie(ofs, cleanup) = get_output_stream(filename, ".pftrace");
auto amount_read = std::atomic<size_t>{0};
auto is_done = std::promise<void>{};
auto _mtx = std::mutex{};
auto _reader = [ofs, &_mtx, &is_done, &amount_read](
::perfetto::TracingSession::ReadTraceCallbackArgs _args) {
auto _lk = std::unique_lock<std::mutex>{_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<void>{};
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();
@@ -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 <deque>
namespace rocprofiler
{
namespace tool
{
void
write_perfetto(
tool_table* tool_functions,
uint64_t pid,
std::vector<rocprofiler_agent_v0_t> agent_data,
std::deque<rocprofiler_buffer_tracing_hip_api_record_t>* hip_api_data,
std::deque<rocprofiler_buffer_tracing_hsa_api_record_t>* hsa_api_data,
std::deque<rocprofiler_buffer_tracing_kernel_dispatch_record_t>* kernel_dispatch_data,
std::deque<rocprofiler_buffer_tracing_memory_copy_record_t>* memory_copy_data,
std::deque<rocprofiler_buffer_tracing_marker_api_record_t>* marker_api_data,
std::deque<rocprofiler_buffer_tracing_scratch_memory_record_t>* scratch_api_data);
} // namespace tool
} // namespace rocprofiler
@@ -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);