AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Update to use rocprofiler-sdk (#55)

Browse Source 
- Renames the CMake option "ROCPROFSYS_USE_HIP" to "ROCPROFSYS_USE_ROCM"
- Remove the "ROCPROFSYS_USE_ROCM_SMI option. Controlled with the "ROCPROFSYS_USE_ROCM" option, instead.
   - Runtime configuration can still toggle ROCPROFSYS_USE_ROCM_SMI to disable the sampling.
- Rename ROCPROFSYS_HIP_VERSION macro to ROCPROFSYS_ROCM_VERSION and remove blocks for `ROCPROFSYS_ROCM_VERSION < 60000`
- Remove ROCPROFSYS_USE_ROCTRACER and ROCPROFSYS_USE_ROCPROFILER
- Update test cases
- Update docker files and workflows to install cmake 3.21, which is required for the rocprofiler-sdk findPackage script.
- Removed rocm-6.2 from workflows due to a rocprofiler-sdk API change.
This commit is contained in:
David Galiffi
2024-12-13 18:48:39 -05:00
committed by GitHub
parent d3725df816
commit 88aa2d3cbe
87 changed files with 3842 additions and 6261 deletions
Download Patch File Download Diff File Expand all files Collapse all files
source/lib/rocprof-sys/library/CMakeLists.txt
+6 -21
View File
@@ -23,8 +23,7 @@ set(library_headers
${CMAKE_CURRENT_LIST_DIR}/rcclp.hpp
${CMAKE_CURRENT_LIST_DIR}/rocm.hpp
${CMAKE_CURRENT_LIST_DIR}/rocm_smi.hpp
${CMAKE_CURRENT_LIST_DIR}/rocprofiler.hpp
${CMAKE_CURRENT_LIST_DIR}/roctracer.hpp
${CMAKE_CURRENT_LIST_DIR}/rocprofiler-sdk.hpp
${CMAKE_CURRENT_LIST_DIR}/runtime.hpp
${CMAKE_CURRENT_LIST_DIR}/sampling.hpp
${CMAKE_CURRENT_LIST_DIR}/thread_data.hpp
@@ -35,37 +34,23 @@ set(library_headers
target_sources(rocprofiler-systems-object-library PRIVATE ${library_sources}
${library_headers})
if(ROCPROFSYS_USE_ROCTRACER OR ROCPROFSYS_USE_ROCPROFILER)
target_sources(rocprofiler-systems-object-library
PRIVATE ${CMAKE_CURRENT_LIST_DIR}/rocm.cpp)
endif()
if(ROCPROFSYS_USE_ROCTRACER)
target_sources(rocprofiler-systems-object-library
PRIVATE ${CMAKE_CURRENT_LIST_DIR}/roctracer.cpp)
endif()
if(ROCPROFSYS_USE_RCCL)
target_sources(rocprofiler-systems-object-library
PRIVATE ${CMAKE_CURRENT_LIST_DIR}/rcclp.cpp)
endif()
if(ROCPROFSYS_USE_ROCPROFILER)
if(ROCPROFSYS_USE_ROCM)
target_sources(
rocprofiler-systems-object-library
PRIVATE ${CMAKE_CURRENT_LIST_DIR}/rocprofiler.cpp
${CMAKE_CURRENT_LIST_DIR}/rocprofiler.hpp)
endif()
if(ROCPROFSYS_USE_ROCM_SMI)
target_sources(rocprofiler-systems-object-library
PRIVATE ${CMAKE_CURRENT_LIST_DIR}/rocm_smi.cpp)
PRIVATE ${CMAKE_CURRENT_LIST_DIR}/rocm.cpp
${CMAKE_CURRENT_LIST_DIR}/rocprofiler-sdk.cpp
${CMAKE_CURRENT_LIST_DIR}/rocm_smi.cpp)
add_subdirectory(rocprofiler-sdk)
endif()
add_subdirectory(causal)
add_subdirectory(components)
add_subdirectory(coverage)
add_subdirectory(rocm)
add_subdirectory(tracing)
set(ndebug_sources
source/lib/rocprof-sys/library/components/CMakeLists.txt
-12
View File
@@ -28,8 +28,6 @@ set(component_headers
${CMAKE_CURRENT_LIST_DIR}/mpi_gotcha.hpp
${CMAKE_CURRENT_LIST_DIR}/numa_gotcha.hpp
${CMAKE_CURRENT_LIST_DIR}/rcclp.hpp
${CMAKE_CURRENT_LIST_DIR}/rocprofiler.hpp
${CMAKE_CURRENT_LIST_DIR}/roctracer.hpp
${CMAKE_CURRENT_LIST_DIR}/pthread_gotcha.hpp
${CMAKE_CURRENT_LIST_DIR}/pthread_create_gotcha.hpp
${CMAKE_CURRENT_LIST_DIR}/pthread_mutex_gotcha.hpp)
@@ -37,16 +35,6 @@ set(component_headers
target_sources(rocprofiler-systems-object-library PRIVATE ${component_sources}
${component_headers})
if(ROCPROFSYS_USE_ROCPROFILER)
target_sources(rocprofiler-systems-object-library
PRIVATE ${CMAKE_CURRENT_LIST_DIR}/rocprofiler.cpp)
endif()
if(ROCPROFSYS_USE_ROCTRACER)
target_sources(rocprofiler-systems-object-library
PRIVATE ${CMAKE_CURRENT_LIST_DIR}/roctracer.cpp)
endif()
if(ROCPROFSYS_USE_RCCL)
target_sources(rocprofiler-systems-object-library
PRIVATE ${CMAKE_CURRENT_LIST_DIR}/rcclp.cpp)
source/lib/rocprof-sys/library/components/category_region.hpp
+4 -4
View File
@@ -64,13 +64,13 @@ using tim::type_list;
// these categories increment push/pop counts, which are used for sanity checks since
// they should ALWAYS be popped if they were pushed
using tracing_count_categories_t =
type_list<category::host, category::mpi, category::pthread, category::rocm_hip,
category::rocm_hsa, category::rocm_rccl>;
type_list<category::host, category::mpi, category::pthread, category::rocm_hip_api,
category::rocm_hsa_api, category::rocm_rccl>;
// convert these categories to throughput points
using causal_throughput_categories_t =
type_list<category::host, category::kokkos, category::ompt, category::rocm_hip,
category::rocm_hsa, category::rocm_rccl, category::rocm_roctx>;
type_list<category::host, category::kokkos, category::ompt, category::rocm_hip_api,
category::rocm_hsa_api, category::rocm_rccl, category::rocm_marker_api>;
// define this outside of category region functions so that the
// static thread_local is global instead of per-template instantiation
source/lib/rocprof-sys/library/components/pthread_create_gotcha.cpp
+1 -5
View File
@@ -28,7 +28,6 @@
#include "core/utility.hpp"
#include "library/causal/delay.hpp"
#include "library/components/category_region.hpp"
#include "library/components/roctracer.hpp"
#include "library/runtime.hpp"
#include "library/sampling.hpp"
#include "library/thread_data.hpp"
@@ -61,7 +60,7 @@ shutdown();
namespace component
{
using bundle_t = tim::lightweight_tuple<comp::wall_clock, comp::roctracer_data>;
using bundle_t = tim::lightweight_tuple<comp::wall_clock>;
using category_region_t = tim::lightweight_tuple<category_region<category::pthread>>;
namespace
@@ -82,7 +81,6 @@ inline void
start_bundle(bundle_t& _bundle, int64_t _tid, Args&&... _args)
{
if(!get_use_timemory() && !get_use_perfetto()) return;
trait::runtime_enabled<comp::roctracer_data>::set(get_use_roctracer());
ROCPROFSYS_BASIC_VERBOSE_F(3, "starting bundle '%s' in thread %li...\n",
_bundle.key().c_str(), _tid);
if constexpr(sizeof...(Args) > 0)
@@ -619,5 +617,3 @@ pthread_create_gotcha::operator()(pthread_t* thread, const pthread_attr_t* attr,
}
} // namespace component
} // namespace rocprofsys
TIMEMORY_INITIALIZE_STORAGE(component::roctracer_data)
source/lib/rocprof-sys/library/components/rocprofiler.cpp
-193
View File
@@ -1,193 +0,0 @@
// MIT License
//
// Copyright (c) 2022-2024 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 "library/components/rocprofiler.hpp"
#include "core/common.hpp"
#include "core/config.hpp"
#include "core/debug.hpp"
#include "core/defines.hpp"
#include "core/dynamic_library.hpp"
#include "core/perfetto.hpp"
#include "core/redirect.hpp"
#include "library/rocprofiler.hpp"
#include "library/sampling.hpp"
#include "library/thread_data.hpp"
#include <timemory/storage/types.hpp>
#include <timemory/utility/types.hpp>
#include <timemory/variadic/functional.hpp>
#include <timemory/variadic/lightweight_tuple.hpp>
#include <rocprofiler.h>
#include <cstdint>
#include <string_view>
#include <type_traits>
namespace rocprofsys
{
namespace component
{
namespace
{
auto&
rocprofiler_activity_count()
{
static std::atomic<int64_t> _v{ 0 };
return _v;
}
} // namespace
unique_ptr_t<rocm_data_t>&
rocm_data(int64_t _tid)
{
using thread_data_t = thread_data<rocm_data_t, rocm_event>;
return thread_data_t::instance(construct_on_thread{ _tid });
}
rocm_event::rocm_event(uint32_t _dev, uint32_t _thr, uint32_t _queue,
std::string _event_name, rocm_metric_type _begin,
rocm_metric_type _end, uint32_t _feature_count, void* _features_v)
: device_id{ _dev }
, thread_id{ _thr }
, queue_id{ _queue }
, entry{ _begin }
, exit{ _end }
, name(std::move(_event_name))
{
feature_values.reserve(_feature_count);
feature_names.reserve(_feature_count);
auto* _features = static_cast<rocprofiler_feature_t*>(_features_v);
for(uint32_t i = 0; i < _feature_count; ++i)
{
const rocprofiler_feature_t* p = &_features[i];
feature_names.emplace_back(i);
switch(p->data.kind)
{
// Output metrics results
case ROCPROFILER_DATA_KIND_UNINIT: break;
case ROCPROFILER_DATA_KIND_BYTES:
feature_values.emplace_back(
rocm_feature_value{ p->data.result_bytes.size });
break;
case ROCPROFILER_DATA_KIND_INT32:
feature_values.emplace_back(rocm_feature_value{ p->data.result_int32 });
break;
case ROCPROFILER_DATA_KIND_FLOAT:
feature_values.emplace_back(rocm_feature_value{ p->data.result_float });
break;
case ROCPROFILER_DATA_KIND_DOUBLE:
feature_values.emplace_back(rocm_feature_value{ p->data.result_double });
break;
case ROCPROFILER_DATA_KIND_INT64:
feature_values.emplace_back(rocm_feature_value{ p->data.result_int64 });
break;
}
}
}
std::string
rocm_event::as_string() const
{
std::stringstream _ss{};
_ss << name << ", device: " << device_id << ", queue: " << queue_id
<< ", thread: " << thread_id << ", entry: " << entry << ", exit = " << exit;
_ss.precision(3);
_ss << std::fixed;
for(size_t i = 0; i < feature_names.size(); ++i)
{
auto _name = rocprofsys::rocprofiler::get_data_labels().at(device_id).at(
feature_names.at(i));
_ss << ", " << _name << " = ";
auto _as_string = [&_ss](auto&& itr) { _ss << std::setw(4) << itr; };
std::visit(_as_string, feature_values.at(i));
}
return _ss.str();
}
void
rocprofiler::preinit()
{
rocprofiler_data::label() = "rocprofiler";
rocprofiler_data::description() = "ROCm hardware counters";
}
void
rocprofiler::start()
{
if(tracker_type::start() == 0) setup();
}
void
rocprofiler::stop()
{
if(tracker_type::stop() == 0) shutdown();
}
bool
rocprofiler::is_setup()
{
return rocprofsys::rocprofiler::is_setup();
}
void
rocprofiler::add_setup(const std::string&, std::function<void()>&&)
{}
void
rocprofiler::add_shutdown(const std::string&, std::function<void()>&&)
{}
void
rocprofiler::remove_setup(const std::string&)
{}
void
rocprofiler::remove_shutdown(const std::string&)
{}
void
rocprofiler::setup()
{
ROCPROFSYS_VERBOSE_F(1, "rocprofiler is setup\n");
}
void
rocprofiler::shutdown()
{
rocprofsys::rocprofiler::post_process();
rocprofsys::rocprofiler::rocm_cleanup();
ROCPROFSYS_VERBOSE_F(1, "rocprofiler is shutdown\n");
}
scope::transient_destructor
rocprofiler::protect_flush_activity()
{
return scope::transient_destructor([]() { --rocprofiler_activity_count(); },
[]() { ++rocprofiler_activity_count(); });
}
} // namespace component
} // namespace rocprofsys
ROCPROFSYS_INSTANTIATE_EXTERN_COMPONENT(rocprofiler, false, void)
ROCPROFSYS_INSTANTIATE_EXTERN_COMPONENT(rocprofiler_data, true,
tim::component::rocprofiler_value)
source/lib/rocprof-sys/library/components/rocprofiler.hpp
-241
View File
@@ -1,241 +0,0 @@
// MIT License
//
// Copyright (c) 2022-2024 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 "core/components/fwd.hpp"
#include "core/defines.hpp"
#include "library/thread_data.hpp"
#include <timemory/api.hpp>
#include <timemory/backends/hardware_counters.hpp>
#include <timemory/components/base.hpp>
#include <timemory/components/data_tracker/components.hpp>
#include <timemory/components/macros.hpp>
#include <timemory/enum.h>
#include <timemory/macros.hpp>
#include <timemory/macros/os.hpp>
#include <timemory/mpl/concepts.hpp>
#include <timemory/mpl/macros.hpp>
#include <timemory/mpl/type_traits.hpp>
#include <timemory/mpl/types.hpp>
#include <timemory/utility/transient_function.hpp>
#include <array>
#include <cstdint>
#include <string>
#include <string_view>
#include <variant>
#include <vector>
namespace rocprofsys
{
namespace component
{
using rocm_metric_type = unsigned long long;
using rocm_info_entry = ::tim::hardware_counters::info;
using rocm_feature_value = std::variant<uint32_t, float, uint64_t, double>;
struct rocm_counter
{
std::array<rocm_metric_type, ROCPROFSYS_ROCM_MAX_COUNTERS> counters;
};
struct rocm_event
{
using value_type = rocm_feature_value;
uint32_t device_id = 0;
uint32_t thread_id = 0;
uint32_t queue_id = 0;
rocm_metric_type entry = 0;
rocm_metric_type exit = 0;
std::string name = {};
std::vector<size_t> feature_names = {};
std::vector<rocm_feature_value> feature_values = {};
rocm_event() = default;
rocm_event(uint32_t _dev, uint32_t _thr, uint32_t _queue, std::string _event_name,
rocm_metric_type begin, rocm_metric_type end, uint32_t _feature_count,
void* _features);
std::string as_string() const;
friend std::ostream& operator<<(std::ostream& _os, const rocm_event& _v)
{
return (_os << _v.as_string());
}
friend bool operator<(const rocm_event& _lhs, const rocm_event& _rhs)
{
return std::tie(_lhs.device_id, _lhs.queue_id, _lhs.entry, _lhs.thread_id) <
std::tie(_rhs.device_id, _rhs.queue_id, _rhs.entry, _rhs.thread_id);
}
};
using rocm_data_t = std::vector<rocm_event>;
using rocm_data_tracker = data_tracker<rocm_feature_value, rocm_event>;
rocprofsys::unique_ptr_t<rocm_data_t>&
rocm_data(int64_t _tid = threading::get_id());
using rocprofiler_value = typename rocm_event::value_type;
using rocprofiler_data = data_tracker<rocprofiler_value, rocprofiler>;
struct rocprofiler
: base<rocprofiler, void>
, private policy::instance_tracker<rocprofiler, false>
{
using value_type = void;
using base_type = base<rocprofiler, void>;
using tracker_type = policy::instance_tracker<rocprofiler, false>;
ROCPROFSYS_DEFAULT_OBJECT(rocprofiler)
static void preinit();
static void global_init() { setup(); }
static void global_finalize() { shutdown(); }
static bool is_setup();
static void setup();
static void shutdown();
static void add_setup(const std::string&, std::function<void()>&&);
static void add_shutdown(const std::string&, std::function<void()>&&);
static void remove_setup(const std::string&);
static void remove_shutdown(const std::string&);
void start();
void stop();
// this function protects rocprofiler_flush_activty from being called
// when rocprof-sys exits during a callback
[[nodiscard]] static scope::transient_destructor protect_flush_activity();
};
#if !defined(ROCPROFSYS_USE_ROCPROFILER)
inline void
rocprofiler::setup()
{}
inline void
rocprofiler::shutdown()
{}
inline bool
rocprofiler::is_setup()
{
return false;
}
#endif
} // namespace component
} // namespace rocprofsys
namespace tim
{
namespace component
{
using ::rocprofsys::component::rocm_data_tracker;
using ::rocprofsys::component::rocm_feature_value;
using ::rocprofsys::component::rocprofiler_data;
using ::rocprofsys::component::rocprofiler_value;
} // namespace component
} // namespace tim
namespace tim
{
namespace operation
{
template <>
struct set_storage<component::rocm_data_tracker>
{
using T = component::rocm_data_tracker;
static constexpr size_t max_threads = 4096;
using type = T;
using storage_array_t = std::array<storage<type>*, max_threads>;
friend struct get_storage<component::rocm_data_tracker>;
ROCPROFSYS_DEFAULT_OBJECT(set_storage)
auto operator()(storage<type>*, size_t) const {}
auto operator()(type&, size_t) const {}
auto operator()(storage<type>* _v) const { get().fill(_v); }
private:
static storage_array_t& get()
{
static storage_array_t _v = { nullptr };
return _v;
}
};
template <>
struct get_storage<component::rocm_data_tracker>
{
using type = component::rocm_data_tracker;
ROCPROFSYS_DEFAULT_OBJECT(get_storage)
auto operator()(const type&) const
{
return operation::set_storage<type>::get().at(0);
}
auto operator()() const
{
type _obj{};
return (*this)(_obj);
}
auto operator()(size_t _idx) const
{
return operation::set_storage<type>::get().at(_idx);
}
auto operator()(type&, size_t _idx) const { return (*this)(_idx); }
};
} // namespace operation
} // namespace tim
#if !defined(ROCPROFSYS_USE_ROCPROFILER)
ROCPROFSYS_DEFINE_CONCRETE_TRAIT(is_available, component::rocprofiler_data, false_type)
#endif
TIMEMORY_SET_COMPONENT_API(component::rocprofiler_data, project::timemory,
category::timing, os::supports_unix)
ROCPROFSYS_DEFINE_CONCRETE_TRAIT(is_timing_category, component::rocprofiler_data,
false_type)
ROCPROFSYS_DEFINE_CONCRETE_TRAIT(uses_timing_units, component::rocprofiler_data,
false_type)
ROCPROFSYS_DEFINE_CONCRETE_TRAIT(report_units, component::rocprofiler_data, false_type)
TIMEMORY_STATISTICS_TYPE(component::rocprofiler_data, component::rocprofiler_value)
TIMEMORY_STATISTICS_TYPE(component::rocm_data_tracker, component::rocm_feature_value)
ROCPROFSYS_DEFINE_CONCRETE_TRAIT(report_units, component::rocm_data_tracker, false_type)
#if !defined(ROCPROFSYS_EXTERN_COMPONENTS) || \
(defined(ROCPROFSYS_EXTERN_COMPONENTS) && ROCPROFSYS_EXTERN_COMPONENTS > 0)
# include <timemory/operations.hpp>
ROCPROFSYS_DECLARE_EXTERN_COMPONENT(rocprofiler, false, void)
ROCPROFSYS_DECLARE_EXTERN_COMPONENT(rocprofiler_data, true, double)
#endif
source/lib/rocprof-sys/library/components/roctracer.cpp
-396
View File
@@ -1,396 +0,0 @@
// MIT License
//
// Copyright (c) 2022-2024 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 "library/components/roctracer.hpp"
#include "core/common.hpp"
#include "core/config.hpp"
#include "core/debug.hpp"
#include "core/defines.hpp"
#include "core/dynamic_library.hpp"
#include "core/redirect.hpp"
#include "library/roctracer.hpp"
#include "library/runtime.hpp"
#include "library/thread_data.hpp"
#include "library/thread_info.hpp"
#include <chrono>
#include <roctracer.h>
#define HIP_PROF_HIP_API_STRING 1
#include <roctracer_ext.h>
#include <roctracer_hip.h>
#if ROCPROFSYS_HIP_VERSION < 50300
# include <roctracer_hcc.h>
#endif
#define AMD_INTERNAL_BUILD 1
#include <roctracer_hsa.h>
namespace rocprofsys
{
namespace component
{
namespace
{
auto&
roctracer_activity_count()
{
static std::atomic<int64_t> _v{ 0 };
return _v;
}
} // namespace
void
roctracer::preinit()
{
roctracer_data::label() = "roctracer";
roctracer_data::description() = "ROCm tracer (activity API)";
}
void
roctracer::start()
{
if(tracker_type::start() == 0) setup(nullptr);
}
void
roctracer::stop()
{
if(tracker_type::stop() == 0) shutdown();
}
bool
roctracer::is_setup()
{
return roctracer_is_setup();
}
void
roctracer::add_setup(const std::string& _lbl, std::function<void()>&& _func)
{
roctracer_setup_routines().emplace_back(_lbl, std::move(_func));
}
void
roctracer::add_shutdown(const std::string& _lbl, std::function<void()>&& _func)
{
roctracer_shutdown_routines().emplace_back(_lbl, std::move(_func));
}
void
roctracer::remove_setup(const std::string& _lbl)
{
auto& _data = roctracer_setup_routines();
for(auto itr = _data.begin(); itr != _data.end(); ++itr)
{
if(itr->first == _lbl)
{
_data.erase(itr);
break;
}
}
}
void
roctracer::remove_shutdown(const std::string& _lbl)
{
auto& _data = roctracer_setup_routines();
for(auto itr = _data.begin(); itr != _data.end(); ++itr)
{
if(itr->first == _lbl)
{
_data.erase(itr);
break;
}
}
}
void
roctracer::setup(void* table, bool on_load_trace)
{
if(!get_use_roctracer()) return;
auto_lock_t _lk{ type_mutex<roctracer>() };
if(roctracer_is_setup()) return;
roctracer_is_setup() = true;
ROCPROFSYS_VERBOSE_F(1, "setting up roctracer...\n");
ROCPROFSYS_SCOPED_SAMPLING_ON_CHILD_THREADS(false);
dynamic_library _amdhip64{ "ROCPROFSYS_ROCTRACER_LIBAMDHIP64",
find_library_path("libamdhip64.so",
{ "ROCPROFSYS_ROCM_PATH", "ROCM_PATH" },
{ ROCPROFSYS_DEFAULT_ROCM_PATH }) };
#if ROCPROFSYS_HIP_VERSION_MAJOR == 4 && ROCPROFSYS_HIP_VERSION_MINOR < 4
dynamic_library _kfdwrapper{
"ROCPROFSYS_ROCTRACER_LIBKFDWRAPPER",
find_library_path("libkfdwrapper64.so", { "ROCPROFSYS_ROCM_PATH", "ROCM_PATH" },
{ ROCPROFSYS_DEFAULT_ROCM_PATH },
{ "roctracer/lib", "roctracer/lib64", "lib", "lib64" })
};
#endif
ROCPROFSYS_ROCTRACER_CALL(roctracer_set_properties(ACTIVITY_DOMAIN_HIP_API, nullptr));
// Allocating tracing pool
roctracer_properties_t properties{};
memset(&properties, 0, sizeof(roctracer_properties_t));
// properties.mode = 0x1000;
properties.buffer_size = 0x100;
properties.buffer_callback_fun = hip_activity_callback;
ROCPROFSYS_ROCTRACER_CALL(roctracer_open_pool(&properties));
#if ROCPROFSYS_HIP_VERSION_MAJOR == 4 && ROCPROFSYS_HIP_VERSION_MINOR >= 4
// HIP 4.5.0 has an invalid warning
redirect _rd{ std::cerr, "roctracer_enable_callback(), get_op_end(), invalid domain "
"ID(4) in: roctracer_enable_callback(hip_api_callback, "
"nullptr)roctracer_enable_activity_expl(), get_op_end(), "
"invalid domain ID(4) in: roctracer_enable_activity()" };
#endif
if(get_trace_hip_api())
{
ROCPROFSYS_ROCTRACER_CALL(roctracer_enable_domain_callback(
ACTIVITY_DOMAIN_HIP_API, hip_api_callback, nullptr));
}
if(get_use_roctx())
{
ROCPROFSYS_ROCTRACER_CALL(roctracer_enable_domain_callback(
ACTIVITY_DOMAIN_ROCTX, roctx_api_callback, nullptr));
}
if(get_trace_hip_activity())
{
// Enable HIP activity tracing
ROCPROFSYS_ROCTRACER_CALL(
roctracer_enable_domain_activity(ACTIVITY_DOMAIN_HIP_OPS));
}
if(table != nullptr)
{
ROCPROFSYS_VERBOSE(1 || on_load_trace, "[OnLoad] setting up HSA...\n");
bool trace_hsa_api = get_trace_hsa_api();
// Enable HSA API callbacks/activity
if(trace_hsa_api)
{
std::vector<std::string> hsa_api_vec =
tim::delimit(get_trace_hsa_api_types());
// initialize HSA tracing
roctracer_set_properties(
static_cast<activity_domain_t>(ACTIVITY_DOMAIN_HSA_API), (void*) table);
if(!hsa_api_vec.empty())
{
for(const auto& itr : hsa_api_vec)
{
uint32_t cid = HSA_API_ID_NUMBER;
const char* api = itr.c_str();
ROCPROFSYS_ROCTRACER_CALL(roctracer_op_code(
static_cast<activity_domain_t>(ACTIVITY_DOMAIN_HSA_API), api,
&cid, nullptr));
ROCPROFSYS_ROCTRACER_CALL(roctracer_enable_op_callback(
static_cast<activity_domain_t>(ACTIVITY_DOMAIN_HSA_API), cid,
hsa_api_callback, nullptr));
ROCPROFSYS_VERBOSE(1 || on_load_trace, " HSA-trace(%s)", api);
}
}
else
{
ROCPROFSYS_VERBOSE(1 || on_load_trace, " HSA-trace()\n");
ROCPROFSYS_ROCTRACER_CALL(roctracer_enable_domain_callback(
static_cast<activity_domain_t>(ACTIVITY_DOMAIN_HSA_API),
hsa_api_callback, nullptr));
}
}
bool trace_hsa_activity = get_trace_hsa_activity();
// Enable HSA GPU activity
if(trace_hsa_activity)
{
#if ROCPROFSYS_HIP_VERSION < 50300
using namespace roctracer;
// initialize HSA tracing
const char* output_prefix = nullptr;
hsa_ops_properties_t ops_properties{
table, reinterpret_cast<activity_async_callback_t>(hsa_activity_callback),
nullptr, output_prefix
};
#elif ROCPROFSYS_HIP_VERSION < 50301
hsa_ops_properties_t ops_properties;
ops_properties.table = table;
ops_properties.reserved1[0] = reinterpret_cast<void*>(&hsa_activity_callback);
ops_properties.reserved1[1] = nullptr;
ops_properties.reserved1[2] = nullptr;
#else
hsa_ops_properties_t ops_properties{
table, reinterpret_cast<void*>(&hsa_activity_callback), nullptr, nullptr
};
#endif
roctracer_set_properties(
static_cast<activity_domain_t>(ACTIVITY_DOMAIN_HSA_OPS), &ops_properties);
ROCPROFSYS_VERBOSE(1 || on_load_trace, " HSA-activity-trace()\n");
ROCPROFSYS_ROCTRACER_CALL(roctracer_enable_op_activity(
static_cast<activity_domain_t>(ACTIVITY_DOMAIN_HSA_OPS), HSA_OP_ID_COPY));
}
}
// callback for HSA
for(auto& itr : roctracer_setup_routines())
itr.second();
// make sure all async callbacks are allocated
for(size_t i = 0; i < thread_info::get_peak_num_threads(); ++i)
hip_exec_activity_callbacks(i);
ROCPROFSYS_VERBOSE_F(1, "roctracer is setup\n");
}
void
roctracer::flush()
{
auto wait_for_activity_flush_completion = []() {
uint16_t nitr = 0;
while(roctracer_activity_count() > 0 && nitr++ < 10)
std::this_thread::sleep_for(std::chrono::milliseconds{ 100 });
};
// a flush may already be happening
wait_for_activity_flush_completion();
if(roctracer_activity_count() == 0)
{
ROCPROFSYS_VERBOSE_F(2, "executing roctracer_flush_activity()...\n");
ROCPROFSYS_ROCTRACER_CALL(roctracer_flush_activity());
// wait to make sure flush completes
std::this_thread::sleep_for(std::chrono::milliseconds{ 100 });
wait_for_activity_flush_completion();
}
else
{
ROCPROFSYS_CI_FAIL(true,
"roctracer_activity_count() != 0 (== %li). "
"roctracer::shutdown() most likely called during abort",
roctracer_activity_count().load());
}
ROCPROFSYS_VERBOSE_F(2, "executing hip_exec_activity_callbacks(0..%zu)\n",
thread_info::get_peak_num_threads());
// make sure all async operations are executed
for(size_t i = 0; i < thread_info::get_peak_num_threads(); ++i)
hip_exec_activity_callbacks(i);
ROCPROFSYS_VERBOSE_F(2, "roctracer flush completed\n");
}
void
roctracer::shutdown()
{
auto_lock_t _lk{ type_mutex<roctracer>() };
if(!roctracer_is_setup()) return;
roctracer_is_setup() = false;
ROCPROFSYS_VERBOSE_F(1, "shutting down roctracer...\n");
// callback for hsa
ROCPROFSYS_VERBOSE_F(2, "executing %zu roctracer_shutdown_routines...\n",
roctracer_shutdown_routines().size());
for(auto& itr : roctracer_shutdown_routines())
itr.second();
#if ROCPROFSYS_HIP_VERSION_MAJOR == 4 && ROCPROFSYS_HIP_VERSION_MINOR >= 4
ROCPROFSYS_DEBUG_F("redirecting roctracer warnings\n");
// HIP 4.5.0 has an invalid warning
redirect _rd{
std::cerr, "roctracer_disable_callback(), get_op_end(), invalid domain ID(4) "
"in: roctracer_disable_callback()roctracer_disable_activity(), "
"get_op_end(), invalid domain ID(4) in: roctracer_disable_activity()"
};
#endif
if(get_trace_hip_api())
{
ROCPROFSYS_VERBOSE_F(
2,
"executing roctracer_disable_domain_callback(ACTIVITY_DOMAIN_HIP_API)...\n");
ROCPROFSYS_ROCTRACER_CALL(
roctracer_disable_domain_callback(ACTIVITY_DOMAIN_HIP_API));
}
if(get_use_roctx())
{
ROCPROFSYS_VERBOSE_F(
2, "executing roctracer_disable_domain_activity(ACTIVITY_DOMAIN_ROCTX)...\n");
ROCPROFSYS_ROCTRACER_CALL(
roctracer_disable_domain_callback(ACTIVITY_DOMAIN_ROCTX));
}
if(get_trace_hip_activity())
{
ROCPROFSYS_VERBOSE_F(
2,
"executing roctracer_disable_domain_activity(ACTIVITY_DOMAIN_HIP_OPS)...\n");
ROCPROFSYS_ROCTRACER_CALL(
roctracer_disable_domain_activity(ACTIVITY_DOMAIN_HIP_OPS));
}
if(get_trace_hsa_api())
{
ROCPROFSYS_VERBOSE_F(
2,
"executing roctracer_disable_domain_activity(ACTIVITY_DOMAIN_HSA_API)...\n");
ROCPROFSYS_ROCTRACER_CALL(
roctracer_disable_domain_callback(ACTIVITY_DOMAIN_HSA_API));
}
if(get_trace_hsa_api())
{
ROCPROFSYS_VERBOSE_F(
2, "executing roctracer_disable_op_activity(ACTIVITY_DOMAIN_HSA_OPS, "
"HSA_OP_ID_COPY)...\n");
ROCPROFSYS_ROCTRACER_CALL(
roctracer_disable_op_activity(ACTIVITY_DOMAIN_HSA_OPS, HSA_OP_ID_COPY));
}
ROCPROFSYS_VERBOSE_F(1, "roctracer is shutdown\n");
}
scope::transient_destructor
roctracer::protect_flush_activity()
{
return scope::transient_destructor([]() { --roctracer_activity_count(); },
[]() { ++roctracer_activity_count(); });
}
} // namespace component
} // namespace rocprofsys
ROCPROFSYS_INSTANTIATE_EXTERN_COMPONENT(roctracer, false, void)
ROCPROFSYS_INSTANTIATE_EXTERN_COMPONENT(roctracer_data, true, double)
source/lib/rocprof-sys/library/components/roctracer.hpp
-117
View File
@@ -1,117 +0,0 @@
// MIT License
//
// Copyright (c) 2022-2024 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 "core/common.hpp"
#include "core/components/fwd.hpp"
#include "core/defines.hpp"
#include <timemory/api.hpp>
#include <timemory/components/base.hpp>
#include <timemory/components/data_tracker/components.hpp>
#include <timemory/components/macros.hpp>
#include <timemory/enum.h>
#include <timemory/macros/os.hpp>
#include <timemory/mpl/type_traits.hpp>
#include <timemory/mpl/types.hpp>
#include <timemory/utility/transient_function.hpp>
ROCPROFSYS_COMPONENT_ALIAS(roctracer_data,
::tim::component::data_tracker<double, roctracer>)
namespace rocprofsys
{
namespace component
{
struct roctracer
: base<roctracer, void>
, private policy::instance_tracker<roctracer, false>
{
using value_type = void;
using base_type = base<roctracer, void>;
using tracker_type = policy::instance_tracker<roctracer, false>;
ROCPROFSYS_DEFAULT_OBJECT(roctracer)
static void preinit();
static void global_finalize() { shutdown(); }
static bool is_setup();
static void setup(void* hsa_api_table, bool on_load_trace = false);
static void flush();
static void shutdown();
static void add_setup(const std::string&, std::function<void()>&&);
static void add_shutdown(const std::string&, std::function<void()>&&);
static void remove_setup(const std::string&);
static void remove_shutdown(const std::string&);
void start();
void stop();
// this function protects roctracer_flush_activty from being called
// when rocprof-sys exits during a callback
[[nodiscard]] static scope::transient_destructor protect_flush_activity();
};
#if !defined(ROCPROFSYS_USE_ROCTRACER)
inline void
roctracer::setup(void*, bool)
{}
inline void
roctracer::flush()
{}
inline void
roctracer::shutdown()
{}
inline bool
roctracer::is_setup()
{
return false;
}
#endif
} // namespace component
} // namespace rocprofsys
#if !defined(ROCPROFSYS_USE_ROCTRACER)
ROCPROFSYS_DEFINE_CONCRETE_TRAIT(is_available, component::roctracer_data, false_type)
#endif
TIMEMORY_SET_COMPONENT_API(rocprofsys::component::roctracer_data, project::timemory,
category::timing, os::supports_unix)
ROCPROFSYS_DEFINE_CONCRETE_TRAIT(is_timing_category, component::roctracer_data, true_type)
ROCPROFSYS_DEFINE_CONCRETE_TRAIT(uses_timing_units, component::roctracer_data, true_type)
#if defined(ROCPROFSYS_USE_ROCTRACER) && ROCPROFSYS_USE_ROCTRACER > 0
# if !defined(ROCPROFSYS_EXTERN_COMPONENTS) || \
(defined(ROCPROFSYS_EXTERN_COMPONENTS) && ROCPROFSYS_EXTERN_COMPONENTS > 0)
# include <timemory/operations.hpp>
ROCPROFSYS_DECLARE_EXTERN_COMPONENT(roctracer, false, void)
ROCPROFSYS_DECLARE_EXTERN_COMPONENT(roctracer_data, true, double)
# endif
#endif
source/lib/rocprof-sys/library/rocm.cpp
+8 -202
View File
@@ -25,12 +25,8 @@
#include "core/debug.hpp"
#include "core/dynamic_library.hpp"
#include "core/gpu.hpp"
#include "library/components/rocprofiler.hpp"
#include "library/components/roctracer.hpp"
#include "library/rocm/hsa_rsrc_factory.hpp"
#include "library/rocm_smi.hpp"
#include "library/rocprofiler.hpp"
#include "library/roctracer.hpp"
#include "library/rocprofiler-sdk.hpp"
#include "library/runtime.hpp"
#include "library/thread_data.hpp"
#include "library/tracing.hpp"
@@ -46,208 +42,18 @@
#include <mutex>
#include <tuple>
#if defined(ROCPROFSYS_USE_ROCPROFILER) && ROCPROFSYS_USE_ROCPROFILER > 0
# include <rocprofiler.h>
#if defined(ROCPROFSYS_USE_ROCM) && ROCPROFSYS_USE_ROCM > 0
# include <rocprofiler-sdk/rocprofiler.h>
#endif
using namespace rocprofsys;
namespace rocprofsys
{
namespace rocm
{
std::mutex rocm_mutex = {};
bool is_loaded = false;
bool on_load_trace = (get_env<int>("ROCP_ONLOAD_TRACE", 0) > 0);
std::vector<hardware_counter_info>
rocm_events()
{
return rocprofiler_sdk::get_rocm_events_info();
}
} // namespace rocm
} // namespace rocprofsys
#if defined(ROCPROFSYS_USE_ROCPROFILER) && ROCPROFSYS_USE_ROCPROFILER > 0
std::ostream&
operator<<(std::ostream& _os, const rocprofiler_settings_t& _v)
{
# define ROCPROF_SETTING_FIELD_STR(NAME) JOIN('=', # NAME, _v.NAME)
_os << JOIN(
", ", ROCPROF_SETTING_FIELD_STR(intercept_mode),
ROCPROF_SETTING_FIELD_STR(code_obj_tracking),
ROCPROF_SETTING_FIELD_STR(memcopy_tracking),
ROCPROF_SETTING_FIELD_STR(trace_size), ROCPROF_SETTING_FIELD_STR(trace_local),
ROCPROF_SETTING_FIELD_STR(timeout), ROCPROF_SETTING_FIELD_STR(timestamp_on),
ROCPROF_SETTING_FIELD_STR(hsa_intercepting),
ROCPROF_SETTING_FIELD_STR(k_concurrent), ROCPROF_SETTING_FIELD_STR(opt_mode),
ROCPROF_SETTING_FIELD_STR(obj_dumping));
return _os;
}
#endif
// HSA-runtime tool on-load method
extern "C"
{
#if defined(ROCPROFSYS_USE_ROCPROFILER) && ROCPROFSYS_USE_ROCPROFILER > 0
void OnUnloadTool()
{
ROCPROFSYS_BASIC_VERBOSE_F(2 || rocm::on_load_trace, "Unloading...\n");
rocm::lock_t _lk{ rocm::rocm_mutex, std::defer_lock };
if(!_lk.owns_lock()) _lk.lock();
if(!rocm::is_loaded)
{
ROCPROFSYS_BASIC_VERBOSE_F(1 || rocm::on_load_trace,
"rocprofiler is not loaded\n");
return;
}
rocm::is_loaded = false;
_lk.unlock();
// stop_top_level_timer_if_necessary();
// Final resources cleanup
rocprofsys::rocprofiler::rocm_cleanup();
}
void OnLoadToolProp(rocprofiler_settings_t* settings)
{
using ::rocprofiler::util::HsaRsrcFactory;
if(!config::get_use_rocprofiler() || config::get_rocm_events().empty()) return;
ROCPROFSYS_BASIC_VERBOSE_F(2 || rocm::on_load_trace, "Loading...\n");
rocm::lock_t _lk{ rocm::rocm_mutex, std::defer_lock };
if(!_lk.owns_lock()) _lk.lock();
if(rocm::is_loaded)
{
ROCPROFSYS_BASIC_VERBOSE_F(1 || rocm::on_load_trace,
"rocprofiler is already loaded\n");
return;
}
rocm::is_loaded = true;
_lk.unlock();
// Enable timestamping
settings->timestamp_on = 1;
settings->intercept_mode = 1;
settings->hsa_intercepting = 1;
settings->k_concurrent = 0;
settings->obj_dumping = 0;
// settings->code_obj_tracking = 0;
// settings->memcopy_tracking = 0;
// settings->trace_local = 1;
// settings->opt_mode = 1;
// settings->trace_size = 0;
// settings->timeout = 0;
ROCPROFSYS_BASIC_VERBOSE_F(1 || rocm::on_load_trace, "rocprofiler settings: %s\n",
JOIN("", *settings).c_str());
// Initialize profiling
rocprofsys::rocprofiler::rocm_initialize();
HsaRsrcFactory::Instance().PrintGpuAgents("ROCm");
}
#endif
bool OnLoad(HsaApiTable* table, uint64_t runtime_version, uint64_t failed_tool_count,
const char* const* failed_tool_names)
{
tim::consume_parameters(table, runtime_version, failed_tool_count,
failed_tool_names);
static bool _once = false;
if(_once) return true;
_once = true;
ROCPROFSYS_BASIC_VERBOSE_F(2 || rocm::on_load_trace, "Loading...\n");
ROCPROFSYS_SCOPED_SAMPLING_ON_CHILD_THREADS(false);
if(!tim::get_env("ROCPROFSYS_INIT_TOOLING", true)) return true;
if(!tim::settings::enabled()) return true;
roctracer_is_init() = true;
ROCPROFSYS_BASIC_VERBOSE_F(1 || rocm::on_load_trace, "Loading ROCm tooling...\n");
if(!config::settings_are_configured() && get_state() < State::Active)
rocprofsys_init_tooling_hidden();
ROCPROFSYS_SCOPED_THREAD_STATE(ThreadState::Internal);
#if ROCPROFSYS_HIP_VERSION < 50300
ROCPROFSYS_VERBOSE_F(1 || rocm::on_load_trace,
"Computing the roctracer clock skew...\n");
(void) rocprofsys::get_clock_skew();
#endif
if(get_use_process_sampling() && get_use_rocm_smi())
{
ROCPROFSYS_VERBOSE_F(1 || rocm::on_load_trace,
"Setting rocm_smi state to active...\n");
rocm_smi::set_state(State::Active);
}
comp::roctracer::setup(static_cast<void*>(table), rocm::on_load_trace);
#if defined(ROCPROFSYS_USE_ROCPROFILER) && ROCPROFSYS_USE_ROCPROFILER > 0
bool _force_rocprofiler_init =
tim::get_env("ROCPROFSYS_FORCE_ROCPROFILER_INIT", false, false);
#else
bool _force_rocprofiler_init = false;
#endif
bool _success = true;
bool _is_empty =
(config::settings_are_configured() && config::get_rocm_events().empty());
if(_force_rocprofiler_init || (get_use_rocprofiler() && !_is_empty))
{
#if ROCPROFSYS_HIP_VERSION < 50500
auto _rocprof = dynamic_library{
"ROCPROFSYS_ROCPROFILER_LIBRARY",
find_library_path(
"librocprofiler64.so", { "ROCPROFSYS_ROCM_PATH", "ROCM_PATH" },
{ ROCPROFSYS_DEFAULT_ROCM_PATH },
{ "lib", "lib64", "rocprofiler/lib", "rocprofiler/lib64" }),
(RTLD_LAZY | RTLD_GLOBAL), false
};
ROCPROFSYS_VERBOSE_F(1 || rocm::on_load_trace,
"Loading rocprofiler library (%s=%s)...\n",
_rocprof.envname.c_str(), _rocprof.filename.c_str());
_rocprof.open();
on_load_t _rocprof_load = nullptr;
_success = _rocprof.invoke("OnLoad", _rocprof_load, table, runtime_version,
failed_tool_count, failed_tool_names);
ROCPROFSYS_CONDITIONAL_PRINT_F(!_success,
"Warning! Invoking rocprofiler's OnLoad "
"failed! ROCPROFSYS_ROCPROFILER_LIBRARY=%s\n",
_rocprof.filename.c_str());
ROCPROFSYS_CI_THROW(!_success,
"Warning! Invoking rocprofiler's OnLoad "
"failed! ROCPROFSYS_ROCPROFILER_LIBRARY=%s\n",
_rocprof.filename.c_str());
#endif
}
else
{
using ::rocprofiler::util::HsaRsrcFactory;
HsaRsrcFactory::Instance().PrintGpuAgents("ROCm");
}
gpu::add_hip_device_metadata();
ROCPROFSYS_BASIC_VERBOSE_F(2 || rocm::on_load_trace, "Loading... %s\n",
(_success) ? "Done" : "Failed");
return _success;
}
// HSA-runtime on-unload method
void OnUnload()
{
ROCPROFSYS_BASIC_VERBOSE_F(2 || rocm::on_load_trace, "Unloading...\n");
rocprofsys_finalize_hidden();
ROCPROFSYS_BASIC_VERBOSE_F(2 || rocm::on_load_trace, "Unloading... Done\n");
}
}
source/lib/rocprof-sys/library/rocm.hpp
+26 -14
View File
@@ -23,36 +23,48 @@
#pragma once
#include "core/defines.hpp"
#include "core/timemory.hpp"
#if defined(ROCPROFSYS_USE_ROCPROFILER) && ROCPROFSYS_USE_ROCPROFILER > 0
# include <rocprofiler.h>
#if defined(ROCPROFSYS_USE_ROCM) && ROCPROFSYS_USE_ROCM > 0
# include <rocprofiler-sdk/registration.h>
# include <rocprofiler-sdk/rocprofiler.h>
#endif
#include <cstdint>
#include <mutex>
#include <vector>
namespace rocprofsys
{
namespace rocm
{
using lock_t = std::unique_lock<std::mutex>;
using hardware_counter_info = ::tim::hardware_counters::info;
extern std::mutex rocm_mutex;
extern bool is_loaded;
std::vector<hardware_counter_info>
rocm_events();
#if !defined(ROCPROFSYS_USE_ROCM) || ROCPROFSYS_USE_ROCM == 0
inline std::vector<hardware_counter_info>
rocm_events()
{
return std::vector<hardware_counter_info>();
}
#endif
} // namespace rocm
} // namespace rocprofsys
extern "C"
{
struct HsaApiTable;
using on_load_t = bool (*)(HsaApiTable*, uint64_t, uint64_t, const char* const*);
struct rocprofiler_tool_configure_result_t;
struct rocprofiler_client_id_t;
bool OnLoad(HsaApiTable* table, uint64_t runtime_version, uint64_t failed_tool_count,
const char* const* failed_tool_names) ROCPROFSYS_PUBLIC_API;
void OnUnload() ROCPROFSYS_PUBLIC_API;
using rocprofiler_configure_t =
rocprofiler_tool_configure_result_t* (*) (uint32_t version,
const char* runtime_version,
uint32_t priority,
rocprofiler_client_id_t* client_id);
#if defined(ROCPROFSYS_USE_ROCPROFILER) && ROCPROFSYS_USE_ROCPROFILER > 0
void OnLoadToolProp(rocprofiler_settings_t* settings) ROCPROFSYS_PUBLIC_API;
void OnUnloadTool() ROCPROFSYS_PUBLIC_API;
#endif
rocprofiler_tool_configure_result_t* rocprofiler_configure(
uint32_t version, const char* runtime_version, uint32_t priority,
rocprofiler_client_id_t* client_id) ROCPROFSYS_PUBLIC_API;
}
source/lib/rocprof-sys/library/rocm/CMakeLists.txt
-7
View File
@@ -1,7 +0,0 @@
#
if(ROCPROFSYS_USE_ROCPROFILER OR ROCPROFSYS_USE_ROCTRACER)
target_sources(
rocprofiler-systems-object-library
PRIVATE ${CMAKE_CURRENT_LIST_DIR}/hsa_rsrc_factory.hpp
${CMAKE_CURRENT_LIST_DIR}/hsa_rsrc_factory.cpp)
endif()
source/lib/rocprof-sys/library/rocm/hsa_rsrc_factory.cpp
-1027
View File
File diff suppressed because it is too large Load Diff
source/lib/rocprof-sys/library/rocm/hsa_rsrc_factory.hpp
-582
View File
@@ -1,582 +0,0 @@
// MIT License
//
// Copyright (c) 2022-2024 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 "core/exception.hpp"
#define AMD_INTERNAL_BUILD 1
#include <hsa.h>
#include <hsa_api_trace.h>
#include <hsa_ext_amd.h>
#include <hsa_ext_finalize.h>
#include <hsa_ven_amd_aqlprofile.h>
#include <hsa_ven_amd_loader.h>
#include <atomic>
#include <cstdint>
#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <ctime>
#include <iostream>
#include <map>
#include <mutex>
#include <string>
#include <vector>
#define HSA_ARGUMENT_ALIGN_BYTES 16
#define HSA_QUEUE_ALIGN_BYTES 64
#define HSA_PACKET_ALIGN_BYTES 64
#define HSA_MESSAGE_LENGTH 4096
#define CHECK_STATUS(msg, status) \
do \
{ \
if((status) != HSA_STATUS_SUCCESS) \
{ \
const char* emsg = 0; \
hsa_status_string(status, &emsg); \
char _buffer[HSA_MESSAGE_LENGTH]; \
snprintf(_buffer, HSA_MESSAGE_LENGTH - 1, "%s: %s", msg, \
emsg ? emsg : "<unknown error>"); \
throw ::rocprofsys::exception<std::runtime_error>(_buffer); \
} \
} while(0)
#define CHECK_ITER_STATUS(msg, status) \
do \
{ \
if((status) != HSA_STATUS_INFO_BREAK) \
{ \
const char* emsg = 0; \
hsa_status_string(status, &emsg); \
char _buffer[HSA_MESSAGE_LENGTH]; \
snprintf(_buffer, HSA_MESSAGE_LENGTH - 1, "%s: %s", msg, \
emsg ? emsg : "<unknown error>"); \
throw ::rocprofsys::exception<std::runtime_error>(_buffer); \
} \
} while(0)
namespace rocprofiler
{
namespace util
{
static const size_t MEM_PAGE_BYTES = 0x1000;
static const size_t MEM_PAGE_MASK = MEM_PAGE_BYTES - 1;
typedef decltype(hsa_agent_t::handle) hsa_agent_handle_t;
struct hsa_pfn_t
{
decltype(::hsa_init)* hsa_init;
decltype(::hsa_shut_down)* hsa_shut_down;
decltype(::hsa_agent_get_info)* hsa_agent_get_info;
decltype(::hsa_iterate_agents)* hsa_iterate_agents;
decltype(::hsa_queue_create)* hsa_queue_create;
decltype(::hsa_queue_destroy)* hsa_queue_destroy;
decltype(::hsa_queue_load_read_index_relaxed)* hsa_queue_load_read_index_relaxed;
decltype(::hsa_queue_load_write_index_relaxed)* hsa_queue_load_write_index_relaxed;
decltype(
::hsa_queue_add_write_index_scacq_screl)* hsa_queue_add_write_index_scacq_screl;
decltype(::hsa_signal_create)* hsa_signal_create;
decltype(::hsa_signal_destroy)* hsa_signal_destroy;
decltype(::hsa_signal_load_relaxed)* hsa_signal_load_relaxed;
decltype(::hsa_signal_store_relaxed)* hsa_signal_store_relaxed;
decltype(::hsa_signal_wait_scacquire)* hsa_signal_wait_scacquire;
decltype(::hsa_signal_store_screlease)* hsa_signal_store_screlease;
decltype(::hsa_code_object_reader_create_from_file)*
hsa_code_object_reader_create_from_file;
decltype(::hsa_executable_create_alt)* hsa_executable_create_alt;
decltype(
::hsa_executable_load_agent_code_object)* hsa_executable_load_agent_code_object;
decltype(::hsa_executable_freeze)* hsa_executable_freeze;
decltype(::hsa_executable_destroy)* hsa_executable_destroy;
decltype(::hsa_executable_get_symbol)* hsa_executable_get_symbol;
decltype(::hsa_executable_symbol_get_info)* hsa_executable_symbol_get_info;
decltype(::hsa_executable_iterate_symbols)* hsa_executable_iterate_symbols;
decltype(::hsa_system_get_info)* hsa_system_get_info;
decltype(
::hsa_system_get_major_extension_table)* hsa_system_get_major_extension_table;
decltype(::hsa_amd_agent_iterate_memory_pools)* hsa_amd_agent_iterate_memory_pools;
decltype(::hsa_amd_memory_pool_get_info)* hsa_amd_memory_pool_get_info;
decltype(::hsa_amd_memory_pool_allocate)* hsa_amd_memory_pool_allocate;
decltype(::hsa_amd_agents_allow_access)* hsa_amd_agents_allow_access;
decltype(::hsa_amd_memory_async_copy)* hsa_amd_memory_async_copy;
decltype(::hsa_amd_signal_async_handler)* hsa_amd_signal_async_handler;
decltype(
::hsa_amd_profiling_set_profiler_enabled)* hsa_amd_profiling_set_profiler_enabled;
decltype(
::hsa_amd_profiling_get_async_copy_time)* hsa_amd_profiling_get_async_copy_time;
decltype(::hsa_amd_profiling_get_dispatch_time)* hsa_amd_profiling_get_dispatch_time;
};
// Encapsulates information about a Hsa Agent such as its
// handle, name, max queue size, max wavefront size, etc.
struct AgentInfo
{
// Handle of Agent
hsa_agent_t dev_id;
// Agent type - Cpu = 0, Gpu = 1 or Dsp = 2
uint32_t dev_type;
// APU flag
bool is_apu;
// Agent system index
uint32_t dev_index;
// GFXIP name
char gfxip[64];
// Name of Agent whose length is less than 64
char name[64];
// Max size of Wavefront size
uint32_t max_wave_size;
// Max size of Queue buffer
uint32_t max_queue_size;
// Hsail profile supported by agent
hsa_profile_t profile;
// CPU/GPU/kern-arg memory pools
hsa_amd_memory_pool_t cpu_pool;
hsa_amd_memory_pool_t gpu_pool;
hsa_amd_memory_pool_t kern_arg_pool;
// The number of compute unit available in the agent.
uint32_t cu_num;
// Maximum number of waves possible in a Compute Unit.
uint32_t waves_per_cu;
// Number of SIMD's per compute unit CU
uint32_t simds_per_cu;
// Number of Shader Engines (SE) in Gpu
uint32_t se_num;
// Number of Shader Arrays Per Shader Engines in Gpu
uint32_t shader_arrays_per_se;
// SGPR/VGPR/LDS block sizes
uint32_t sgpr_block_dflt;
uint32_t sgpr_block_size;
uint32_t vgpr_block_size;
static const uint32_t lds_block_size = 128 * 4;
};
// HSA timer class
// Provides current HSA timestampa and system-clock/ns conversion API
class HsaTimer
{
public:
typedef uint64_t timestamp_t;
static const timestamp_t TIMESTAMP_MAX = UINT64_MAX;
typedef long double freq_t;
enum time_id_t
{
TIME_ID_CLOCK_REALTIME = 0,
TIME_ID_CLOCK_REALTIME_COARSE = 1,
TIME_ID_CLOCK_MONOTONIC = 2,
TIME_ID_CLOCK_MONOTONIC_COARSE = 3,
TIME_ID_CLOCK_MONOTONIC_RAW = 4,
TIME_ID_NUMBER
};
HsaTimer(const hsa_pfn_t* hsa_api)
: hsa_api_(hsa_api)
{
timestamp_t sysclock_hz = 0;
hsa_status_t status = hsa_api_->hsa_system_get_info(
HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY, &sysclock_hz);
CHECK_STATUS("hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY)", status);
sysclock_factor_ = (freq_t) 1000000000 / (freq_t) sysclock_hz;
}
// Methods for system-clock/ns conversion
timestamp_t sysclock_to_ns(const timestamp_t& sysclock) const
{
return timestamp_t((freq_t) sysclock * sysclock_factor_);
}
timestamp_t ns_to_sysclock(const timestamp_t& time) const
{
return timestamp_t((freq_t) time / sysclock_factor_);
}
// Method for timespec/ns conversion
static timestamp_t timespec_to_ns(const timespec& time)
{
return ((timestamp_t) time.tv_sec * 1000000000) + time.tv_nsec;
}
// Return timestamp in 'ns'
timestamp_t timestamp_ns() const
{
timestamp_t sysclock;
hsa_status_t status =
hsa_api_->hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP, &sysclock);
CHECK_STATUS("hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP)", status);
return sysclock_to_ns(sysclock);
}
// Return time in 'ns'
timestamp_t clocktime_ns(clockid_t clock_id) const
{
timespec time;
clock_gettime(clock_id, &time);
return timespec_to_ns(time);
}
// Return pair of correlated values of profiling timestamp and time with
// correlation error for a given time ID and number of iterations
void correlated_pair_ns(time_id_t time_id, uint32_t iters, timestamp_t* timestamp_v,
timestamp_t* time_v, timestamp_t* error_v)
{
clockid_t clock_id = 0;
switch(time_id)
{
case TIME_ID_CLOCK_REALTIME: clock_id = CLOCK_REALTIME; break;
case TIME_ID_CLOCK_REALTIME_COARSE: clock_id = CLOCK_REALTIME_COARSE; break;
case TIME_ID_CLOCK_MONOTONIC: clock_id = CLOCK_MONOTONIC; break;
case TIME_ID_CLOCK_MONOTONIC_COARSE: clock_id = CLOCK_MONOTONIC_COARSE; break;
case TIME_ID_CLOCK_MONOTONIC_RAW: clock_id = CLOCK_MONOTONIC_RAW; break;
default: CHECK_STATUS("internal error: invalid time_id", HSA_STATUS_ERROR);
}
std::vector<timestamp_t> ts_vec(iters);
std::vector<timespec> tm_vec(iters);
const uint32_t steps = iters - 1;
for(uint32_t i = 0; i < iters; ++i)
{
hsa_api_->hsa_system_get_info(HSA_SYSTEM_INFO_TIMESTAMP, &ts_vec[i]);
clock_gettime(clock_id, &tm_vec[i]);
}
const timestamp_t ts_base = sysclock_to_ns(ts_vec.front());
const timestamp_t tm_base = timespec_to_ns(tm_vec.front());
const timestamp_t error = (ts_vec.back() - ts_vec.front()) / (2 * steps);
timestamp_t ts_accum = 0;
timestamp_t tm_accum = 0;
for(uint32_t i = 0; i < iters; ++i)
{
ts_accum += (ts_vec[i] - ts_base);
tm_accum += (timespec_to_ns(tm_vec[i]) - tm_base);
}
*timestamp_v = (ts_accum / iters) + ts_base + error;
*time_v = (tm_accum / iters) + tm_base;
*error_v = error;
}
private:
// Timestamp frequency factor
freq_t sysclock_factor_;
// HSA API table
const hsa_pfn_t* const hsa_api_;
};
class HsaRsrcFactory
{
public:
static const size_t CMD_SLOT_SIZE_B = 0x40;
typedef std::recursive_mutex mutex_t;
typedef HsaTimer::timestamp_t timestamp_t;
static HsaRsrcFactory* Create(bool initialize_hsa = true)
{
std::lock_guard<mutex_t> lck(mutex_);
HsaRsrcFactory* obj = instance_.load(std::memory_order_relaxed);
if(obj == nullptr)
{
obj = new HsaRsrcFactory(initialize_hsa);
instance_.store(obj, std::memory_order_release);
}
return obj;
}
static HsaRsrcFactory& Instance()
{
HsaRsrcFactory* obj = instance_.load(std::memory_order_acquire);
if(obj == nullptr) obj = Create(false);
hsa_status_t status = (obj != nullptr) ? HSA_STATUS_SUCCESS : HSA_STATUS_ERROR;
CHECK_STATUS("HsaRsrcFactory::Instance() failed", status);
return *obj;
}
static void Destroy()
{
std::lock_guard<mutex_t> lck(mutex_);
if(instance_) delete instance_.load();
instance_ = nullptr;
}
// Return system agent info
const AgentInfo* GetAgentInfo(const hsa_agent_t agent);
// Get the count of Hsa Gpu Agents available on the platform
// @return uint32_t Number of Gpu agents on platform
uint32_t GetCountOfGpuAgents();
// Get the count of Hsa Cpu Agents available on the platform
// @return uint32_t Number of Cpu agents on platform
uint32_t GetCountOfCpuAgents();
// Get the AgentInfo handle of a Gpu device
// @param idx Gpu Agent at specified index
// @param agent_info Output parameter updated with AgentInfo
// @return bool true if successful, false otherwise
bool GetGpuAgentInfo(uint32_t idx, const AgentInfo** agent_info);
// Get the AgentInfo handle of a Cpu device
// @param idx Cpu Agent at specified index
// @param agent_info Output parameter updated with AgentInfo
// @return bool true if successful, false otherwise
bool GetCpuAgentInfo(uint32_t idx, const AgentInfo** agent_info);
// Create a Queue object and return its handle. The queue object is expected
// to support user requested number of Aql dispatch packets.
// @param agent_info Gpu Agent on which to create a queue object
// @param num_Pkts Number of packets to be held by queue
// @param queue Output parameter updated with handle of queue object
// @return bool true if successful, false otherwise
bool CreateQueue(const AgentInfo* agent_info, uint32_t num_pkts, hsa_queue_t** queue);
// Create a Signal object and return its handle.
// @param value Initial value of signal object
// @param signal Output parameter updated with handle of signal object
// @return bool true if successful, false otherwise
bool CreateSignal(uint32_t value, hsa_signal_t* signal);
// Allocate local GPU memory
// @param agent_info Agent from whose memory region to allocate
// @param size Size of memory in terms of bytes
// @return uint8_t* Pointer to buffer, null if allocation fails.
uint8_t* AllocateLocalMemory(const AgentInfo* agent_info, size_t size);
// Allocate memory tp pass kernel parameters
// Memory is alocated accessible for all CPU agents and for GPU given by AgentInfo
// parameter.
// @param agent_info Agent from whose memory region to allocate
// @param size Size of memory in terms of bytes
// @return uint8_t* Pointer to buffer, null if allocation fails.
uint8_t* AllocateKernArgMemory(const AgentInfo* agent_info, size_t size);
// Allocate system memory accessible from both CPU and GPU
// Memory is alocated accessible to all CPU agents and AgentInfo parameter is ignored.
// @param agent_info Agent from whose memory region to allocate
// @param size Size of memory in terms of bytes
// @return uint8_t* Pointer to buffer, null if allocation fails.
uint8_t* AllocateSysMemory(const AgentInfo* agent_info, size_t size);
// Allocate memory for command buffer.
// @param agent_info Agent from whose memory region to allocate
// @param size Size of memory in terms of bytes
// @return uint8_t* Pointer to buffer, null if allocation fails.
uint8_t* AllocateCmdMemory(const AgentInfo* agent_info, size_t size);
// Wait signal
hsa_signal_value_t SignalWait(const hsa_signal_t& signal,
const hsa_signal_value_t& signal_value) const;
// Wait signal with signal value restore
void SignalWaitRestore(const hsa_signal_t& signal,
const hsa_signal_value_t& signal_value) const;
// Copy data from GPU to host memory
bool Memcpy(const hsa_agent_t& agent, void* dst, const void* src, size_t size);
bool Memcpy(const AgentInfo* agent_info, void* dst, const void* src, size_t size);
// Memory free method
static bool FreeMemory(void* ptr);
// Loads an Assembled Brig file and Finalizes it into Device Isa
// @param agent_info Gpu device for which to finalize
// @param brig_path File path of the Assembled Brig file
// @param kernel_name Name of the kernel to finalize
// @param code_desc Handle of finalized Code Descriptor that could
// be used to submit for execution
// @return true if successful, false otherwise
bool LoadAndFinalize(const AgentInfo* agent_info, const char* brig_path,
const char* kernel_name, hsa_executable_t* hsa_exec,
hsa_executable_symbol_t* code_desc);
// Print the various fields of Hsa Gpu Agents
bool PrintGpuAgents(const std::string& header);
// Utils for submitting AQL packet to a given queue
static void* GetSlotPointer(hsa_queue_t* queue, const uint64_t& idx);
static void* GetReadPointer(hsa_queue_t* queue);
static uint64_t Submit(hsa_queue_t* queue, const void* packet);
static uint64_t Submit(hsa_queue_t* queue, const void* packet, size_t size_bytes);
// Enable executables loading tracking
static bool IsExecutableTracking() { return executable_tracking_on_; }
static void EnableExecutableTracking(HsaApiTable* table);
static const char* GetKernelNameRef(uint64_t addr);
// Initialize HSA API table
void static InitHsaApiTable(HsaApiTable* table);
static const hsa_pfn_t* HsaApi() { return &hsa_api_; }
// Return AqlProfile API table
typedef hsa_ven_amd_aqlprofile_pfn_t aqlprofile_pfn_t;
const aqlprofile_pfn_t* AqlProfileApi() const { return &aqlprofile_api_; }
// Return Loader API table
const hsa_ven_amd_loader_1_00_pfn_t* LoaderApi() const { return &loader_api_; }
// Methods for system-clock/ns conversion and timestamp in 'ns'
timestamp_t SysclockToNs(const timestamp_t& sysclock) const
{
return timer_->sysclock_to_ns(sysclock);
}
timestamp_t NsToSysclock(const timestamp_t& time) const
{
return timer_->ns_to_sysclock(time);
}
timestamp_t TimestampNs() const { return timer_->timestamp_ns(); }
timestamp_t GetSysTimeout() const { return timeout_; }
static timestamp_t GetTimeoutNs() { return timeout_ns_; }
static void SetTimeoutNs(const timestamp_t& time)
{
std::lock_guard<mutex_t> lck(mutex_);
timeout_ns_ = time;
if(instance_ != nullptr)
Instance().timeout_ = Instance().timer_->ns_to_sysclock(time);
}
void CorrelateTime(HsaTimer::time_id_t time_id, uint32_t iters)
{
timestamp_t timestamp_v = 0;
timestamp_t time_v = 0;
timestamp_t error_v = 0;
timer_->correlated_pair_ns(time_id, iters, &timestamp_v, &time_v, &error_v);
time_shift_[time_id] = time_v - timestamp_v;
time_error_[time_id] = error_v;
}
hsa_status_t GetTimeVal(uint32_t time_id, uint64_t time_stamp, uint64_t* time_value)
{
if(time_id >= HsaTimer::TIME_ID_NUMBER) return HSA_STATUS_ERROR;
*time_value = time_stamp + time_shift_[time_id];
return HSA_STATUS_SUCCESS;
}
hsa_status_t GetTimeErr(uint32_t time_id, uint64_t* err)
{
*err = time_error_[time_id];
return HSA_STATUS_SUCCESS;
}
private:
// System agents iterating callback
static hsa_status_t GetHsaAgentsCallback(hsa_agent_t agent, void* data);
// Callback function to find and bind kernarg region of an agent
static hsa_status_t FindMemRegionsCallback(hsa_region_t region, void* data);
// Load AQL profile HSA extension library directly
static hsa_status_t LoadAqlProfileLib(aqlprofile_pfn_t* api);
// Constructor of the class. Will initialize the Hsa Runtime and
// query the system topology to get the list of Cpu and Gpu devices
explicit HsaRsrcFactory(bool initialize_hsa);
// Destructor of the class
~HsaRsrcFactory();
// Add an instance of AgentInfo representing a Hsa Gpu agent
const AgentInfo* AddAgentInfo(const hsa_agent_t agent);
// To mmap command buffer memory
static const bool CMD_MEMORY_MMAP = false;
// HSA was initialized
const bool initialize_hsa_;
static std::atomic<HsaRsrcFactory*> instance_;
static mutex_t mutex_;
// Used to maintain a list of Hsa Gpu Agent Info
std::vector<const AgentInfo*> gpu_list_;
std::vector<hsa_agent_t> gpu_agents_;
// Used to maintain a list of Hsa Cpu Agent Info
std::vector<const AgentInfo*> cpu_list_;
std::vector<hsa_agent_t> cpu_agents_;
// System agents map
std::map<hsa_agent_handle_t, const AgentInfo*> agent_map_;
// Executables loading tracking
typedef std::map<uint64_t, const char*> symbols_map_t;
static symbols_map_t* symbols_map_;
static bool executable_tracking_on_;
static void* to_dump_code_obj_;
static hsa_status_t hsa_executable_freeze_interceptor(hsa_executable_t executable,
const char* options);
static hsa_status_t hsa_executable_destroy_interceptor(hsa_executable_t executable);
static hsa_status_t executable_symbols_cb(hsa_executable_t exec,
hsa_executable_symbol_t symbol, void* data);
// HSA runtime API table
static hsa_pfn_t hsa_api_;
// AqlProfile API table
aqlprofile_pfn_t aqlprofile_api_;
// Loader API table
hsa_ven_amd_loader_1_00_pfn_t loader_api_;
// System timeout, ns
static timestamp_t timeout_ns_;
// System timeout, sysclock
timestamp_t timeout_;
// HSA timer
HsaTimer* timer_;
// Time shift array to support time conversion
timestamp_t time_shift_[HsaTimer::TIME_ID_NUMBER];
timestamp_t time_error_[HsaTimer::TIME_ID_NUMBER];
// CPU/kern-arg memory pools
hsa_amd_memory_pool_t* cpu_pool_;
hsa_amd_memory_pool_t* kern_arg_pool_;
};
} // namespace util
} // namespace rocprofiler
source/lib/rocprof-sys/library/rocm_smi.hpp
+3 -2
View File
@@ -128,7 +128,8 @@ private:
static bool shutdown();
};
#if !defined(ROCPROFSYS_USE_ROCM_SMI)
#if !defined(ROCPROFSYS_USE_ROCM) || ROCPROFSYS_USE_ROCM == 0
inline void
setup()
{}
@@ -154,7 +155,7 @@ inline void set_state(State) {}
} // namespace rocm_smi
} // namespace rocprofsys
#if defined(ROCPROFSYS_USE_ROCM_SMI) && ROCPROFSYS_USE_ROCM_SMI > 0
#if defined(ROCPROFSYS_USE_ROCM) && ROCPROFSYS_USE_ROCM > 0
# if !defined(ROCPROFSYS_EXTERN_COMPONENTS) || \
(defined(ROCPROFSYS_EXTERN_COMPONENTS) && ROCPROFSYS_EXTERN_COMPONENTS > 0)
source/lib/rocprof-sys/library/rocprofiler-sdk.cpp
+1308
View File
File diff suppressed because it is too large Load Diff
source/lib/rocprof-sys/library/rocprofiler.hpp → source/lib/rocprof-sys/library/rocprofiler-sdk.hpp
+17 -45
View File
@@ -1,6 +1,6 @@
// MIT License
//
// Copyright (c) 2022-2024 Advanced Micro Devices, Inc. All Rights Reserved.
// Copyright (c) 2022 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
@@ -22,67 +22,39 @@
#pragma once
#include "core/defines.hpp"
#include "core/timemory.hpp"
#include "library/components/rocprofiler.hpp"
#include <timemory/backends/hardware_counters.hpp>
#include <timemory/macros.hpp>
#include <timemory/mpl/concepts.hpp>
#include <timemory/mpl/macros.hpp>
#include <array>
#include <atomic>
#include <cstring>
#include <dlfcn.h>
#include <iostream>
#include <list>
#include <map>
#include <string>
#include <string_view>
#include <tuple>
#include <unistd.h>
#include <utility>
#include <variant>
#include <memory>
#include <vector>
namespace rocprofsys
{
namespace rocprofiler
namespace rocprofiler_sdk
{
std::map<uint32_t, std::vector<std::string_view>>
get_data_labels();
using hardware_counter_info = ::tim::hardware_counters::info;
void
rocm_initialize();
setup();
void
rocm_cleanup();
shutdown();
bool&
is_setup();
void
config();
void
post_process();
std::vector<component::rocm_info_entry>
rocm_metrics();
void
sample();
#if !defined(ROCPROFSYS_USE_ROCPROFILER) || ROCPROFSYS_USE_ROCPROFILER == 0
inline void
post_process()
{}
void
start();
inline void
rocm_cleanup()
{}
void
stop();
inline std::vector<component::rocm_info_entry>
rocm_metrics()
{
return std::vector<component::rocm_info_entry>{};
}
#endif
} // namespace rocprofiler
std::vector<hardware_counter_info>
get_rocm_events_info();
} // namespace rocprofiler_sdk
} // namespace rocprofsys
source/lib/rocprof-sys/library/rocprofiler-sdk/CMakeLists.txt
+9
View File
@@ -0,0 +1,9 @@
#
set(rocprofiler_sdk_sources ${CMAKE_CURRENT_LIST_DIR}/counters.cpp
${CMAKE_CURRENT_LIST_DIR}/fwd.cpp)
set(rocprofiler_sdk_headers ${CMAKE_CURRENT_LIST_DIR}/counters.hpp
${CMAKE_CURRENT_LIST_DIR}/fwd.hpp)
target_sources(rocprofiler-systems-object-library PRIVATE ${rocprofiler_sdk_sources}
${rocprofiler_sdk_headers})
source/lib/rocprof-sys/library/rocprofiler-sdk/counters.cpp
+135
View File
@@ -0,0 +1,135 @@
// MIT License
//
// Copyright (c) 2024 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 "library/rocprofiler-sdk/counters.hpp"
#include "common/synchronized.hpp"
#include "core/debug.hpp"
#include "core/timemory.hpp"
#include "library/rocprofiler-sdk/fwd.hpp"
#include <timemory/utility/types.hpp>
#include <rocprofiler-sdk/agent.h>
#include <rocprofiler-sdk/buffer_tracing.h>
#include <rocprofiler-sdk/callback_tracing.h>
#include <rocprofiler-sdk/cxx/hash.hpp>
#include <rocprofiler-sdk/cxx/name_info.hpp>
#include <rocprofiler-sdk/cxx/operators.hpp>
#include <rocprofiler-sdk/dispatch_counting_service.h>
#include <rocprofiler-sdk/fwd.h>
#include <rocprofiler-sdk/registration.h>
#include <memory>
#include <unordered_map>
#include <vector>
namespace rocprofsys
{
namespace rocprofiler_sdk
{
namespace
{
std::string
get_counter_description(const client_data* tool_data, std::string_view _v)
{
const auto& _info = tool_data->events_info;
for(const auto& itr : _info)
{
if(itr.symbol().find(_v) == 0 || itr.short_description().find(_v) == 0)
{
return itr.long_description();
}
}
return std::string{};
}
} // namespace
void
counter_event::operator()(const client_data* tool_data, ::perfetto::CounterTrack* _track,
timing_interval _timing, scope::config _scope) const
{
if(!record.dispatch_data) return;
const auto& _dispatch_info = record.dispatch_data->dispatch_info;
const auto* _kern_sym_data =
tool_data->get_kernel_symbol_info(_dispatch_info.kernel_id);
auto _bundle = counter_bundle_t{ tim::demangle(_kern_sym_data->kernel_name), _scope };
_bundle.push(_dispatch_info.queue_id.handle)
.start()
.store(record.record_counter.counter_value);
_bundle.stop().pop(_dispatch_info.queue_id.handle);
if(_track && _timing.start > 0 && _timing.end > _timing.start)
{
TRACE_COUNTER(trait::name<category::rocm_counter_collection>::value, *_track,
_timing.start, record.record_counter.counter_value);
TRACE_COUNTER(trait::name<category::rocm_counter_collection>::value, *_track,
_timing.end, 0);
}
}
counter_storage::counter_storage(const client_data* _tool_data, uint64_t _devid,
size_t _idx, std::string_view _name)
: tool_data{ _tool_data }
, device_id{ _devid }
, index{ static_cast<int64_t>(_idx) }
, metric_name{ _name }
, metric_description{ get_counter_description(_tool_data, metric_name) }
{
auto _metric_name = std::string{ _name };
_metric_name =
std::regex_replace(_metric_name, std::regex{ "(.*)\\[([0-9]+)\\]" }, "$1_$2");
storage_name = JOIN('-', "rocprof", "device", device_id, _metric_name);
storage = std::make_unique<counter_storage_type>(tim::standalone_storage{}, index,
storage_name);
{
constexpr auto _unit = ::perfetto::CounterTrack::Unit::UNIT_COUNT;
track_name = JOIN(" ", "GPU", _metric_name, JOIN("", '[', device_id, ']'));
track = std::make_unique<counter_track_type>(
::perfetto::StaticString(track_name.c_str()));
track->set_is_incremental(false);
track->set_unit(_unit);
track->set_unit_multiplier(1);
}
}
void
counter_storage::operator()(const counter_event& _event, timing_interval _timing,
scope::config _scope) const
{
operation::set_storage<counter_data_tracker>{}(storage.get());
_event(tool_data, track.get(), _timing, _scope);
}
void
counter_storage::write() const
{
operation::set_storage<counter_data_tracker>{}(storage.get());
counter_data_tracker::label() = metric_name;
counter_data_tracker::description() = metric_description;
storage->write();
}
} // namespace rocprofiler_sdk
} // namespace rocprofsys
source/lib/rocprof-sys/library/rocprofiler-sdk/counters.hpp
+168
View File
@@ -0,0 +1,168 @@
// MIT License
//
// Copyright (c) 2024 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 "common/synchronized.hpp"
#include "core/debug.hpp"
#include "core/perfetto.hpp"
#include "core/timemory.hpp"
#include "library/rocprofiler-sdk/fwd.hpp"
#include <timemory/utility/types.hpp>
#include <rocprofiler-sdk/agent.h>
#include <rocprofiler-sdk/buffer_tracing.h>
#include <rocprofiler-sdk/callback_tracing.h>
#include <rocprofiler-sdk/cxx/hash.hpp>
#include <rocprofiler-sdk/cxx/name_info.hpp>
#include <rocprofiler-sdk/cxx/operators.hpp>
#include <rocprofiler-sdk/dispatch_counting_service.h>
#include <rocprofiler-sdk/fwd.h>
#include <rocprofiler-sdk/registration.h>
#include <memory>
#include <unordered_map>
#include <vector>
namespace rocprofsys
{
namespace rocprofiler_sdk
{
struct counter_dispatch_record
{
const rocprofiler_dispatch_counting_service_data_t* dispatch_data = nullptr;
rocprofiler_dispatch_id_t dispatch_id = 0;
rocprofiler_counter_id_t counter_id = {};
rocprofiler_record_counter_t record_counter = {};
};
struct counter_data_tag
{};
using counter_data_tracker = component::data_tracker<double, counter_data_tag>;
using counter_storage_type = typename counter_data_tracker::storage_type;
using counter_bundle_t = tim::lightweight_tuple<counter_data_tracker>;
using counter_track_type = ::perfetto::CounterTrack;
struct counter_event
{
ROCPROFSYS_DEFAULT_OBJECT(counter_event)
explicit counter_event(counter_dispatch_record&& _v)
: record{ _v }
{}
void operator()(const client_data* tool_data, counter_track_type*,
timing_interval _timing, scope::config _scope) const;
counter_dispatch_record record = {};
};
struct counter_storage
{
const client_data* tool_data = nullptr;
uint64_t device_id = 0;
int64_t index = 0;
std::string metric_name = {};
std::string metric_description = {};
std::string storage_name = {};
std::string track_name = {};
std::unique_ptr<counter_storage_type> storage = {};
std::unique_ptr<counter_track_type> track = {};
counter_storage(const client_data* _tool_data, uint64_t _devid, size_t _idx,
std::string_view _name);
~counter_storage() = default;
counter_storage(const counter_storage&) = delete;
counter_storage(counter_storage&&) = default;
counter_storage& operator=(const counter_storage&) = delete;
counter_storage& operator=(counter_storage&&) = default;
friend bool operator<(const counter_storage& lhs, const counter_storage& rhs)
{
return std::tie(lhs.storage_name, lhs.device_id, lhs.index) <
std::tie(rhs.storage_name, rhs.device_id, rhs.index);
}
void operator()(const counter_event& _event, timing_interval _timing,
scope::config _scope = scope::get_default()) const;
void write() const;
};
} // namespace rocprofiler_sdk
} // namespace rocprofsys
namespace tim
{
namespace operation
{
template <>
struct set_storage<::rocprofsys::rocprofiler_sdk::counter_data_tracker>
{
static constexpr size_t max_threads = 4096;
using type = ::rocprofsys::rocprofiler_sdk::counter_data_tracker;
using storage_array_t = std::array<storage<type>*, max_threads>;
friend struct get_storage<rocprofsys::rocprofiler_sdk::counter_data_tracker>;
ROCPROFSYS_DEFAULT_OBJECT(set_storage)
auto operator()(storage<type>* _v, size_t _idx) const { get().at(_idx) = _v; }
auto operator()(type&, size_t) const {}
auto operator()(storage<type>* _v) const { get().fill(_v); }
private:
static storage_array_t& get()
{
static storage_array_t _v = { nullptr };
return _v;
}
};
template <>
struct get_storage<::rocprofsys::rocprofiler_sdk::counter_data_tracker>
{
using type = ::rocprofsys::rocprofiler_sdk::counter_data_tracker;
ROCPROFSYS_DEFAULT_OBJECT(get_storage)
auto operator()(const type&) const
{
return operation::set_storage<type>::get().at(0);
}
auto operator()() const
{
type _obj{};
return (*this)(_obj);
}
auto operator()(size_t _idx) const
{
return operation::set_storage<type>::get().at(_idx);
}
auto operator()(type&, size_t _idx) const { return (*this)(_idx); }
};
} // namespace operation
} // namespace tim
source/lib/rocprof-sys/library/rocprofiler-sdk/fwd.cpp
+270
View File
@@ -0,0 +1,270 @@
// MIT License
//
// Copyright (c) 2024 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 "library/rocprofiler-sdk/fwd.hpp"
#include "core/debug.hpp"
#include "core/state.hpp"
#include <timemory/utility/join.hpp>
#include <exception>
#include <rocprofiler-sdk/agent.h>
#include <rocprofiler-sdk/cxx/name_info.hpp>
#include <rocprofiler-sdk/fwd.h>
#include <rocprofiler-sdk/rocprofiler.h>
#include <algorithm>
#include <utility>
namespace rocprofsys
{
namespace rocprofiler_sdk
{
namespace
{
using tool_agent_vec_t = std::vector<tool_agent>;
rocprofiler_status_t
dimensions_info_callback(rocprofiler_counter_id_t /*id*/,
const rocprofiler_record_dimension_info_t* dim_info,
long unsigned int num_dims, void* user_data)
{
auto* dimensions_info =
static_cast<std::vector<rocprofiler_record_dimension_info_t>*>(user_data);
dimensions_info->reserve(num_dims);
for(size_t j = 0; j < num_dims; j++)
dimensions_info->emplace_back(dim_info[j]);
return ROCPROFILER_STATUS_SUCCESS;
}
rocprofiler_status_t
counters_supported_callback(rocprofiler_agent_id_t agent_id,
rocprofiler_counter_id_t* counters, size_t num_counters,
void* user_data)
{
using value_type = typename agent_counter_info_map_t::mapped_type;
auto* data_v = static_cast<agent_counter_info_map_t*>(user_data);
data_v->emplace(agent_id, value_type{});
for(size_t i = 0; i < num_counters; ++i)
{
auto _info = rocprofiler_counter_info_v0_t{};
auto _dim_info = std::vector<rocprofiler_record_dimension_info_t>{};
ROCPROFILER_CALL(rocprofiler_query_counter_info(
counters[i], ROCPROFILER_COUNTER_INFO_VERSION_0, &_info));
// populate local vector
ROCPROFILER_CALL(rocprofiler_iterate_counter_dimensions(
counters[i], dimensions_info_callback, &_dim_info));
if(!_info.is_constant)
data_v->at(agent_id).emplace_back(agent_id, _info, std::move(_dim_info));
}
return ROCPROFILER_STATUS_SUCCESS;
}
agent_counter_info_map_t
get_agent_counter_info(const tool_agent_vec_t& _agents)
{
auto _data = agent_counter_info_map_t{};
for(auto itr : _agents)
{
ROCPROFILER_CALL(rocprofiler_iterate_agent_supported_counters(
itr.agent->id, counters_supported_callback, &_data));
std::sort(_data.at(itr.agent->id).begin(), _data.at(itr.agent->id).end(),
[](const auto& lhs, const auto& rhs) {
return (lhs.id.handle < rhs.id.handle);
});
for(auto& citr : _data.at(itr.agent->id))
{
std::sort(citr.dimension_info.begin(), citr.dimension_info.end(),
[](const auto& lhs, const auto& rhs) { return (lhs.id < rhs.id); });
}
}
return _data;
}
} // namespace
rocprofiler_tool_counter_info_t::rocprofiler_tool_counter_info_t(
rocprofiler_agent_id_t _agent_id, parent_type _info, dimension_info_vec_t&& _dim_info)
: parent_type{ _info }
, agent_id{ _agent_id }
, dimension_info{ std::move(_dim_info) }
{}
void
client_data::initialize()
{
buffered_tracing_info = rocprofiler::sdk::get_buffer_tracing_names();
callback_tracing_info = rocprofiler::sdk::get_callback_tracing_names();
static constexpr auto supported_agent_info_version = ROCPROFILER_AGENT_INFO_VERSION_0;
rocprofiler_query_available_agents_cb_t iterate_cb =
[](rocprofiler_agent_version_t version, const void** agents_arr,
size_t num_agents, void* user_data) {
ROCPROFSYS_CONDITIONAL_ABORT(version != supported_agent_info_version,
"rocprofiler agent info version != expected "
"agent info version (=%i). value: %i\n",
supported_agent_info_version, version);
auto _agents_v = std::vector<rocprofiler_agent_v0_t>{};
for(size_t i = 0; i < num_agents; ++i)
{
const auto* _agent =
static_cast<const rocprofiler_agent_v0_t*>(agents_arr[i]);
_agents_v.emplace_back(*_agent);
}
auto* tool_data_v = as_client_data(user_data);
tool_data_v->set_agents(std::move(_agents_v));
return ROCPROFILER_STATUS_SUCCESS;
};
ROCPROFILER_CALL(rocprofiler_query_available_agents(
supported_agent_info_version, iterate_cb, sizeof(rocprofiler_agent_t), this));
}
void
client_data::initialize_event_info()
{
if(agents.empty()) initialize();
if(agent_counter_info.size() != gpu_agents.size())
agent_counter_info = get_agent_counter_info(gpu_agents);
try
{
using qualifier_t = tim::hardware_counters::qualifier;
using qualifier_vec_t = std::vector<qualifier_t>;
for(const auto& aitr : gpu_agents)
{
auto _dev_index = aitr.device_id;
auto _device_qualifier_sym = JOIN("", ":device=", _dev_index);
auto _device_qualifier =
tim::hardware_counters::qualifier{ true, static_cast<int>(_dev_index),
_device_qualifier_sym,
JOIN(" ", "Device", _dev_index) };
auto _counter_info = agent_counter_info.at(aitr.agent->id);
std::sort(_counter_info.begin(), _counter_info.end(),
[](const rocprofiler_tool_counter_info_t& lhs,
const rocprofiler_tool_counter_info_t& rhs) {
if(lhs.is_constant && rhs.is_constant)
return lhs.id < rhs.id;
else if(lhs.is_constant)
return true;
else if(rhs.is_constant)
return false;
if(!lhs.is_derived && !rhs.is_derived)
return lhs.id < rhs.id;
else if(!lhs.is_derived)
return true;
else if(!rhs.is_derived)
return false;
return lhs.id < rhs.id;
});
for(const auto& ditr : _counter_info)
{
auto _long_desc = std::string{ ditr.description };
auto _units = std::string{};
auto _pysym = std::string{};
if(ditr.is_constant)
{
continue;
}
else if(ditr.is_derived)
{
auto _sym = JOIN("", ditr.name, _device_qualifier_sym);
auto _short_desc = JOIN("", "Derived counter: ", ditr.expression);
events_info.emplace_back(hardware_counter_info(
true, tim::hardware_counters::api::rocm, events_info.size(), 0,
_sym, _pysym, _short_desc, _long_desc, _units,
qualifier_vec_t{ _device_qualifier }));
}
else
{
auto _dim_info = std::vector<std::string>{};
for(const auto& itr : ditr.dimension_info)
{
auto _info = (itr.instance_size > 1)
? JOIN("", itr.name, "[", 0, ":",
itr.instance_size - 1, "]")
: std::string{};
if(!_info.empty()) _dim_info.emplace_back(_info);
}
auto _sym = JOIN("", ditr.name, _device_qualifier_sym);
auto _short_desc = JOIN("", ditr.name, " on device ", _dev_index);
if(!_dim_info.empty())
{
namespace join = ::timemory::join;
_short_desc += JOIN(
"", ". ",
join::join(join::array_config{ ", ", "", "" }, _dim_info));
}
events_info.emplace_back(hardware_counter_info(
true, tim::hardware_counters::api::rocm, events_info.size(), 0,
_sym, _pysym, _short_desc, _long_desc, _units,
qualifier_vec_t{ _device_qualifier }));
}
}
}
} catch(std::exception& _e)
{
ROCPROFSYS_WARNING_F(1, "Constructing ROCm event info failed: %s\n", _e.what());
}
}
void
client_data::set_agents(agent_vec_t&& _agents_v)
{
agents = std::move(_agents_v);
std::sort(agents.begin(), agents.end(),
[](const auto& lhs, const auto& rhs) { return lhs.node_id < rhs.node_id; });
cpu_agents.clear();
gpu_agents.clear();
for(const auto& itr : agents)
{
if(itr.type == ROCPROFILER_AGENT_TYPE_CPU)
cpu_agents.emplace_back(tool_agent{ cpu_agents.size(), &itr });
else if(itr.type == ROCPROFILER_AGENT_TYPE_GPU)
gpu_agents.emplace_back(tool_agent{ gpu_agents.size(), &itr });
}
}
} // namespace rocprofiler_sdk
} // namespace rocprofsys
source/lib/rocprof-sys/library/rocprofiler-sdk/fwd.hpp
+252
View File
@@ -0,0 +1,252 @@
// MIT License
//
// Copyright (c) 2024 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 "common/synchronized.hpp"
#include "core/timemory.hpp"
#include <rocprofiler-sdk/agent.h>
#include <rocprofiler-sdk/buffer_tracing.h>
#include <rocprofiler-sdk/callback_tracing.h>
#include <rocprofiler-sdk/cxx/hash.hpp>
#include <rocprofiler-sdk/cxx/name_info.hpp>
#include <rocprofiler-sdk/cxx/operators.hpp>
#include <rocprofiler-sdk/fwd.h>
#include <rocprofiler-sdk/registration.h>
#include <memory>
#include <vector>
namespace rocprofsys
{
namespace rocprofiler_sdk
{
using hardware_counter_info = ::tim::hardware_counters::info;
using kernel_symbol_data_t =
rocprofiler_callback_tracing_code_object_kernel_symbol_register_data_t;
using kernel_symbol_map_t =
std::unordered_map<rocprofiler_kernel_id_t, kernel_symbol_data_t>;
using callback_arg_array_t = std::vector<std::pair<std::string, std::string>>;
struct code_object_callback_record_t
{
uint64_t timestamp = 0;
rocprofiler_callback_tracing_record_t record = {};
rocprofiler_callback_tracing_code_object_load_data_t payload = {};
};
struct kernel_symbol_callback_record_t
{
uint64_t timestamp = 0;
rocprofiler_callback_tracing_record_t record = {};
kernel_symbol_data_t payload = {};
};
struct rocprofiler_tool_counter_info_t : rocprofiler_counter_info_v0_t
{
using this_type = rocprofiler_tool_counter_info_t;
using parent_type = rocprofiler_counter_info_v0_t;
using dimension_info_vec_t = std::vector<rocprofiler_record_dimension_info_t>;
rocprofiler_tool_counter_info_t(rocprofiler_agent_id_t _agent_id, parent_type _info,
dimension_info_vec_t&& _dim_info);
rocprofiler_tool_counter_info_t() = default;
~rocprofiler_tool_counter_info_t() = default;
rocprofiler_tool_counter_info_t(const rocprofiler_tool_counter_info_t&) = default;
rocprofiler_tool_counter_info_t(rocprofiler_tool_counter_info_t&&) noexcept = default;
rocprofiler_tool_counter_info_t& operator=(const rocprofiler_tool_counter_info_t&) =
default;
rocprofiler_tool_counter_info_t& operator =(
rocprofiler_tool_counter_info_t&&) noexcept = default;
rocprofiler_agent_id_t agent_id = {};
std::vector<rocprofiler_record_dimension_info_t> dimension_info = {};
};
struct tool_agent
{
uint64_t device_id = 0;
const rocprofiler_agent_v0_t* agent = nullptr;
};
struct timing_interval
{
rocprofiler_timestamp_t start = 0;
rocprofiler_timestamp_t end = 0;
};
using agent_counter_info_map_t =
std::unordered_map<rocprofiler_agent_id_t,
std::vector<rocprofiler_tool_counter_info_t>>;
using agent_counter_profile_map_t =
std::unordered_map<rocprofiler_agent_id_t,
std::optional<rocprofiler_profile_config_id_t>>;
using counter_id_vec_t = std::vector<rocprofiler_counter_id_t>;
using agent_counter_id_map_t =
std::unordered_map<rocprofiler_agent_id_t, counter_id_vec_t>;
using backtrace_operation_map_t =
std::unordered_map<rocprofiler_callback_tracing_kind_t,
std::unordered_set<rocprofiler_tracing_operation_t>>;
struct client_data
{
static constexpr size_t num_buffers = 3;
static constexpr size_t num_contexts = 2;
using buffer_name_info_t = rocprofiler::sdk::buffer_name_info_t<std::string_view>;
using callback_name_info_t = rocprofiler::sdk::callback_name_info_t<std::string_view>;
using kernel_symbol_vec_t = std::vector<kernel_symbol_callback_record_t*>;
using code_object_vec_t = std::vector<code_object_callback_record_t>;
using buffer_id_vec_t = std::array<rocprofiler_buffer_id_t, num_buffers>;
using context_id_vec_t = std::array<rocprofiler_context_id_t, num_contexts>;
using agent_vec_t = std::vector<rocprofiler_agent_v0_t>;
rocprofiler_client_id_t* client_id = nullptr;
rocprofiler_client_finalize_t client_fini = nullptr;
rocprofiler_context_id_t primary_ctx = { 0 };
rocprofiler_context_id_t counter_ctx = { 0 };
rocprofiler_buffer_id_t kernel_dispatch_buffer = { 0 };
rocprofiler_buffer_id_t memory_copy_buffer = { 0 };
rocprofiler_buffer_id_t counter_collection_buffer = { 0 };
std::vector<rocprofiler_agent_v0_t> agents = {};
std::vector<tool_agent> cpu_agents = {};
std::vector<tool_agent> gpu_agents = {};
std::vector<hardware_counter_info> events_info = {};
agent_counter_id_map_t agent_events = {};
agent_counter_info_map_t agent_counter_info = {};
agent_counter_profile_map_t agent_counter_profiles = {};
common::synchronized<code_object_vec_t> code_object_records = {};
common::synchronized<kernel_symbol_vec_t> kernel_symbol_records = {};
buffer_name_info_t buffered_tracing_info = {};
callback_name_info_t callback_tracing_info = {};
backtrace_operation_map_t backtrace_operations = {};
void initialize();
void initialize_event_info();
void set_agents(agent_vec_t&& agents);
context_id_vec_t get_contexts() const;
buffer_id_vec_t get_buffers() const;
const rocprofiler_agent_t* get_agent(rocprofiler_agent_id_t _id) const;
const tool_agent* get_gpu_tool_agent(rocprofiler_agent_id_t id) const;
const kernel_symbol_data_t* get_kernel_symbol_info(uint64_t _kernel_id) const;
const rocprofiler_tool_counter_info_t* get_tool_counter_info(
rocprofiler_agent_id_t _agent_id, rocprofiler_counter_id_t _counter_id) const;
};
inline client_data::context_id_vec_t
client_data::get_contexts() const
{
return context_id_vec_t{
primary_ctx,
counter_ctx,
};
}
inline client_data::buffer_id_vec_t
client_data::get_buffers() const
{
return buffer_id_vec_t{
kernel_dispatch_buffer,
memory_copy_buffer,
counter_collection_buffer,
};
}
inline const rocprofiler_agent_t*
client_data::get_agent(rocprofiler_agent_id_t _id) const
{
for(const auto& itr : agents)
if(itr.id == _id) return &itr;
return nullptr;
}
inline const tool_agent*
client_data::get_gpu_tool_agent(rocprofiler_agent_id_t id) const
{
for(const auto& itr : gpu_agents)
if(id == itr.agent->id) return &itr;
return nullptr;
}
inline const kernel_symbol_data_t*
client_data::get_kernel_symbol_info(uint64_t _kernel_id) const
{
return kernel_symbol_records.rlock(
[_kernel_id](const auto& _data) -> const kernel_symbol_data_t* {
for(const auto& itr : _data)
{
if(_kernel_id == itr->payload.kernel_id)
{
return &itr->payload;
break;
}
}
return nullptr;
});
}
inline const rocprofiler_tool_counter_info_t*
client_data::get_tool_counter_info(rocprofiler_agent_id_t _agent_id,
rocprofiler_counter_id_t _counter_id) const
{
for(const auto& itr : agent_counter_info.at(_agent_id))
{
if(itr.id == _counter_id) return &itr;
}
return nullptr;
}
inline constexpr client_data*
as_client_data(void* _ptr)
{
return static_cast<client_data*>(_ptr);
}
} // namespace rocprofiler_sdk
} // namespace rocprofsys
#if !defined(ROCPROFILER_CALL)
# define ROCPROFILER_CALL(result) \
{ \
rocprofiler_status_t ROCPROFSYS_VARIABLE(_rocp_status_, __LINE__) = \
(result); \
if(ROCPROFSYS_VARIABLE(_rocp_status_, __LINE__) != \
ROCPROFILER_STATUS_SUCCESS) \
{ \
auto msg = std::stringstream{}; \
std::string status_msg = rocprofiler_get_status_string( \
ROCPROFSYS_VARIABLE(_rocp_status_, __LINE__)); \
msg << "[" #result "][" << __FILE__ << ":" << __LINE__ << "] " \
<< "rocprofiler-sdk call [" << #result \
<< "] failed with error code " \
<< ROCPROFSYS_VARIABLE(_rocp_status_, __LINE__) \
<< " :: " << status_msg; \
ROCPROFSYS_WARNING(0, "%s\n", msg.str().c_str()); \
} \
}
#endif
source/lib/rocprof-sys/library/rocprofiler.cpp
-834
View File
@@ -1,834 +0,0 @@
// MIT License
//
// Copyright (c) 2022-2024 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 "library/rocprofiler.hpp"
#include "core/common.hpp"
#include "core/config.hpp"
#include "core/debug.hpp"
#include "core/gpu.hpp"
#include "core/perfetto.hpp"
#include "library/rocm.hpp"
#include "library/rocm/hsa_rsrc_factory.hpp"
#include <timemory/backends/hardware_counters.hpp>
#include <timemory/manager.hpp>
#include <timemory/mpl/concepts.hpp>
#include <timemory/storage/types.hpp>
#include <timemory/utility/types.hpp>
#include <rocprofiler.h>
#include <atomic>
#include <cstdlib>
#include <dlfcn.h>
#include <hsa.h>
#include <iostream>
#include <mutex>
#include <sstream>
#include <string.h>
#include <string_view>
#include <type_traits>
#include <unistd.h>
#include <vector>
namespace rocprofsys
{
namespace rocprofiler
{
namespace
{
using ::rocprofiler::util::AgentInfo;
using ::rocprofiler::util::HsaRsrcFactory;
auto&
get_event_names()
{
static auto _v = std::map<uint32_t, std::vector<rocprofiler_feature_t>>{};
return _v;
}
} // namespace
// Error handler
void
fatal(const std::string& msg)
{
ROCPROFSYS_PRINT_F("\n");
ROCPROFSYS_PRINT_F("%s\n", msg.c_str());
abort();
}
// Check returned HSA API status
const char*
rocm_error_string(hsa_status_t _status)
{
const char* _err_string = nullptr;
if(_status != HSA_STATUS_SUCCESS) rocprofiler_error_string(&_err_string);
return _err_string;
}
// Check returned HSA API status
bool
rocm_check_status(hsa_status_t _status, const std::set<hsa_status_t>& _nonfatal = {})
{
if(_status != HSA_STATUS_SUCCESS)
{
if(_nonfatal.count(_status) == 0)
fatal(JOIN(" :: ", "ERROR", rocm_error_string(_status)));
ROCPROFSYS_PRINT_F("Warning! %s\n", rocm_error_string(_status));
return false;
}
return true;
}
// Context stored entry type
struct context_entry_t
{
bool valid;
hsa_agent_t agent;
rocprofiler_group_t group;
rocprofiler_callback_data_t data;
};
// Context callback arg
struct callbacks_arg_t
{
rocprofiler_pool_t** pools;
};
// Handler callback arg
struct handler_arg_t
{
rocprofiler_feature_t* features;
unsigned feature_count;
};
bool&
is_setup()
{
static bool _v = false;
return _v;
}
std::map<uint32_t, std::vector<std::string_view>>
get_data_labels()
{
auto _v = std::map<uint32_t, std::vector<std::string_view>>{};
for(const auto& itr : get_event_names())
{
_v[itr.first] = {};
for(auto vitr : itr.second)
_v[itr.first].emplace_back(std::string_view{ vitr.name });
}
return _v;
}
// Dump stored context entry
void
rocm_dump_context_entry(context_entry_t* entry, rocprofiler_feature_t* features,
unsigned feature_count)
{
volatile std::atomic<bool>* valid =
reinterpret_cast<std::atomic<bool>*>(&entry->valid);
while(valid->load() == false)
sched_yield();
const rocprofiler_dispatch_record_t* record = entry->data.record;
if(!record) return; // there is nothing to do here.
auto _queue_id = entry->data.queue_id;
auto _thread_id = entry->data.thread_id;
auto _dev_id = HsaRsrcFactory::Instance().GetAgentInfo(entry->agent)->dev_index;
auto _kernel_name = std::string{ entry->data.kernel_name };
auto _pos = _kernel_name.find_last_of(')');
if(_pos != std::string::npos) _kernel_name = _kernel_name.substr(0, _pos + 1);
rocprofiler_group_t& group = entry->group;
if(group.context == nullptr)
{
fatal("context is nullptr\n");
}
if(feature_count > 0)
{
rocm_check_status(rocprofiler_group_get_data(&group));
rocm_check_status(rocprofiler_get_metrics(group.context));
}
auto _evt =
component::rocm_event{ _dev_id, _thread_id, _queue_id, _kernel_name,
record->begin, record->end, feature_count, features };
component::rocm_data()->emplace_back(_evt);
}
// Profiling completion handler
// Dump and delete the context entry
// Return true if the context was dumped successfully
bool
rocm_context_handler(const rocprofiler_pool_entry_t* entry, void* arg)
{
// Context entry
context_entry_t* ctx_entry = reinterpret_cast<context_entry_t*>(entry->payload);
handler_arg_t* handler_arg = reinterpret_cast<handler_arg_t*>(arg);
// rocm::lock_t _lk{ rocm::rocm_mutex, std::defer_lock };
// if(!_lk.owns_lock()) _lk.lock();
rocm_dump_context_entry(ctx_entry, handler_arg->features, handler_arg->feature_count);
return true;
}
// Kernel disoatch callback
hsa_status_t
rocm_dispatch_callback(const rocprofiler_callback_data_t* callback_data, void* arg,
rocprofiler_group_t* group)
{
// Passed tool data
hsa_agent_t agent = callback_data->agent;
// Open profiling context
const unsigned gpu_id = HsaRsrcFactory::Instance().GetAgentInfo(agent)->dev_index;
callbacks_arg_t* callbacks_arg = reinterpret_cast<callbacks_arg_t*>(arg);
rocprofiler_pool_t* pool = callbacks_arg->pools[gpu_id];
rocprofiler_pool_entry_t pool_entry{};
rocm_check_status(rocprofiler_pool_fetch(pool, &pool_entry));
// Profiling context entry
rocprofiler_t* context = pool_entry.context;
context_entry_t* entry = reinterpret_cast<context_entry_t*>(pool_entry.payload);
// Get group[0]
rocm_check_status(rocprofiler_get_group(context, 0, group));
// Fill profiling context entry
entry->agent = agent;
entry->group = *group;
entry->data = *callback_data;
entry->data.kernel_name = strdup(callback_data->kernel_name);
reinterpret_cast<std::atomic<bool>*>(&entry->valid)->store(true);
return HSA_STATUS_SUCCESS;
}
unsigned
metrics_input(unsigned _device, rocprofiler_feature_t** ret)
{
// Profiling feature objects
auto _events = tim::delimit(config::get_rocm_events(), ", ;\t\n");
std::vector<std::string> _features = {};
auto _this_device = JOIN("", ":device=", _device);
for(auto itr : _events)
{
ROCPROFSYS_VERBOSE_F(3, "Processing feature '%s' for device %u...\n", itr.c_str(),
_device);
auto _pos = itr.find(":device=");
if(_pos != std::string::npos)
{
if(itr.find(_this_device) != std::string::npos)
{
_features.emplace_back(itr.substr(0, _pos));
}
}
else
{
_features.emplace_back(itr);
}
}
const unsigned feature_count = _features.size();
rocprofiler_feature_t* features = new rocprofiler_feature_t[feature_count];
memset(features, 0, feature_count * sizeof(rocprofiler_feature_t));
// PMC events
for(unsigned i = 0; i < feature_count; ++i)
{
ROCPROFSYS_VERBOSE_F(3, "Adding feature '%s' for device %u...\n",
_features.at(i).c_str(), _device);
features[i].kind = ROCPROFILER_FEATURE_KIND_METRIC;
features[i].name = strdup(_features.at(i).c_str());
features[i].parameters = nullptr;
features[i].parameter_count = 0;
}
*ret = features;
return feature_count;
}
using info_data = std::vector<component::rocm_info_entry>;
hsa_status_t
info_data_callback(const rocprofiler_info_data_t info, void* arg)
{
using qualifier_t = tim::hardware_counters::qualifier;
using qualifier_vec_t = std::vector<qualifier_t>;
auto* _data = static_cast<info_data*>(arg);
auto _dev_index = info.agent_index;
switch(info.kind)
{
case ROCPROFILER_INFO_KIND_METRIC:
{
auto _device_qualifier_sym = JOIN("", ":device=", _dev_index);
auto _device_qualifier =
tim::hardware_counters::qualifier{ true, static_cast<int>(_dev_index),
_device_qualifier_sym,
JOIN(" ", "Device", _dev_index) };
auto _long_desc = std::string{ info.metric.description };
auto _units = std::string{};
auto _pysym = std::string{};
if(info.metric.expr != nullptr)
{
auto _sym = JOIN("", info.metric.name, _device_qualifier_sym);
auto _short_desc = JOIN("", "Derived counter: ", info.metric.expr);
_data->emplace_back(component::rocm_info_entry(
true, tim::hardware_counters::api::rocm, _data->size(), 0, _sym,
_pysym, _short_desc, _long_desc, _units,
qualifier_vec_t{ _device_qualifier }));
}
else
{
if(info.metric.instances == 1)
{
auto _sym = JOIN("", info.metric.name, _device_qualifier_sym);
auto _short_desc =
JOIN("", info.metric.name, " on device ", _dev_index);
_data->emplace_back(component::rocm_info_entry(
true, tim::hardware_counters::api::rocm, _data->size(), 0, _sym,
_pysym, _short_desc, _long_desc, _units,
qualifier_vec_t{ _device_qualifier }));
}
else
{
for(uint32_t i = 0; i < info.metric.instances; ++i)
{
auto _instance_qualifier_sym = JOIN("", '[', i, ']');
auto _instance_qualifier =
tim::hardware_counters::qualifier{ true, static_cast<int>(i),
_instance_qualifier_sym,
JOIN(" ", "Instance", i) };
auto _sym = JOIN("", info.metric.name, _instance_qualifier_sym,
_device_qualifier_sym);
auto _short_desc = JOIN("", info.metric.name, " instance ", i,
" on device ", _dev_index);
_data->emplace_back(component::rocm_info_entry(
true, tim::hardware_counters::api::rocm, _data->size(), 0,
_sym, _pysym, _short_desc, _long_desc, _units,
qualifier_vec_t{ _device_qualifier, _instance_qualifier }));
}
}
}
break;
}
default: printf("wrong info kind %u\n", info.kind); return HSA_STATUS_ERROR;
}
return HSA_STATUS_SUCCESS;
}
std::vector<component::rocm_info_entry>
rocm_metrics()
{
std::vector<component::rocm_info_entry> _data = {};
try
{
(void) HsaRsrcFactory::Instance();
} catch(std::runtime_error& _e)
{
ROCPROFSYS_VERBOSE_F(0, "%s\n", _e.what());
return _data;
}
// Available GPU agents
const unsigned gpu_count = HsaRsrcFactory::Instance().GetCountOfGpuAgents();
std::vector<AgentInfo*> _gpu_agents(gpu_count, nullptr);
for(unsigned i = 0; i < gpu_count; ++i)
{
const AgentInfo* _agent = _gpu_agents[i];
const AgentInfo** _agent_p = &_agent;
HsaRsrcFactory::Instance().GetGpuAgentInfo(i, _agent_p);
if(!rocm_check_status(rocprofiler_iterate_info(
&_agent->dev_id, ROCPROFILER_INFO_KIND_METRIC,
info_data_callback, reinterpret_cast<void*>(&_data)),
{ HSA_STATUS_ERROR_NOT_INITIALIZED }))
{
ROCPROFSYS_WARNING_F(-1, "rocprofiler_iterate_info failed for gpu agent %u\n",
i);
}
}
if(gpu_count > 0 && _data.empty())
{
if(!rocm_check_status(rocprofiler_iterate_info(
nullptr, ROCPROFILER_INFO_KIND_METRIC,
info_data_callback, reinterpret_cast<void*>(&_data)),
{ HSA_STATUS_ERROR_NOT_INITIALIZED }))
{
ROCPROFSYS_WARNING_F(
-1, "rocprofiler_iterate_info failed for %i gpu agents\n", gpu_count);
}
}
auto _settings = tim::settings::shared_instance();
if(_settings)
{
auto ritr = _settings->find("ROCPROFSYS_ROCM_EVENTS");
if(ritr != _settings->end())
{
auto _rocm_events = ritr->second;
if(_rocm_events->get_choices().empty())
{
std::vector<std::string> _choices = {};
_choices.reserve(_data.size());
for(auto itr : _data)
{
if(!itr.symbol().empty()) _choices.emplace_back(itr.symbol());
}
_rocm_events->set_choices(_choices);
}
}
}
return _data;
}
void
rocm_initialize()
{
// Available GPU agents
const unsigned gpu_count = HsaRsrcFactory::Instance().GetCountOfGpuAgents();
(void) rocm_metrics();
// Adding dispatch observer
callbacks_arg_t* callbacks_arg = new callbacks_arg_t{};
callbacks_arg->pools = new rocprofiler_pool_t*[gpu_count];
for(unsigned gpu_id = 0; gpu_id < gpu_count; gpu_id++)
{
// Getting profiling features
rocprofiler_feature_t* features = nullptr;
unsigned feature_count = metrics_input(gpu_id, &features);
if(features)
{
get_event_names()[gpu_id].clear();
get_event_names()[gpu_id].reserve(feature_count);
for(unsigned i = 0; i < feature_count; ++i)
get_event_names().at(gpu_id).emplace_back(features[i]);
}
// Handler arg
handler_arg_t* handler_arg = new handler_arg_t{};
handler_arg->features = features;
handler_arg->feature_count = feature_count;
// Context properties
rocprofiler_pool_properties_t properties{};
properties.num_entries = 100;
properties.payload_bytes = sizeof(context_entry_t);
properties.handler = rocm_context_handler;
properties.handler_arg = handler_arg;
// Getting GPU device info
const AgentInfo* agent_info = nullptr;
if(HsaRsrcFactory::Instance().GetGpuAgentInfo(gpu_id, &agent_info) == false)
{
fprintf(stderr, "GetGpuAgentInfo failed\n");
abort();
}
// Open profiling pool
rocprofiler_pool_t* pool = nullptr;
uint32_t mode = 0; // ROCPROFILER_MODE_SINGLEGROUP
rocm_check_status(rocprofiler_pool_open(agent_info->dev_id, features,
feature_count, &pool, mode, &properties));
callbacks_arg->pools[gpu_id] = pool;
}
rocprofiler_queue_callbacks_t callbacks_ptrs{};
callbacks_ptrs.dispatch = rocm_dispatch_callback;
int err = rocprofiler_set_queue_callbacks(callbacks_ptrs, callbacks_arg);
ROCPROFSYS_VERBOSE_F(3, "err=%d, rocprofiler_set_queue_callbacks\n", err);
is_setup() = true;
}
void
rocm_cleanup()
{
// Unregister dispatch callback
rocm_check_status(rocprofiler_remove_queue_callbacks());
// close profiling pool
// rocm_check_status(rocprofiler_pool_flush(pool));
// rocm_check_status(rocprofiler_pool_close(pool));
}
namespace
{
using rocm_event = component::rocm_event;
using rocm_data_t = component::rocm_data_t;
using rocm_metric_type = component::rocm_metric_type;
using rocm_feature_value = component::rocm_feature_value;
using rocm_data_tracker = component::rocm_data_tracker;
void
post_process_perfetto()
{
using counter_track = perfetto_counter_track<rocm_event>;
static bool _once = false;
if(_once) return;
auto _data = rocm_data_t{};
auto _device_data = std::map<uint32_t, std::vector<rocm_event*>>{};
auto _device_fields = std::map<uint32_t, std::vector<std::string_view>>{};
auto _device_range = std::map<uint32_t, std::set<rocm_metric_type>>{};
for(size_t i = 0; i < ROCPROFSYS_MAX_THREADS; ++i)
{
auto& _v = component::rocm_data(i);
if(_v)
{
_data.reserve(_data.size() + _v->size());
for(auto& itr : *_v)
_data.emplace_back(itr);
}
}
if(_data.empty()) return;
_once = true;
std::sort(_data.begin(), _data.end());
auto _get_events = [](std::vector<rocm_event*>& _inp, rocm_metric_type _ts) {
auto _v = std::vector<rocm_event*>{};
for(const auto& itr : _inp)
{
if(_ts >= itr->entry && _ts <= itr->exit) _v.emplace_back(itr);
if(_ts > itr->exit) break;
}
return _v;
};
{
auto _device_time = std::map<uint32_t, std::set<rocm_metric_type>>{};
for(auto& itr : _data)
{
_device_data[itr.device_id].emplace_back(&itr);
_device_time[itr.device_id].emplace(itr.entry);
_device_time[itr.device_id].emplace(itr.exit);
auto _dev_id = itr.device_id;
if(get_use_perfetto() && !counter_track::exists(_dev_id))
{
auto addendum = [&](auto&& _v) {
return JOIN(" ", "Device", _v, JOIN("", '[', _dev_id, ']'));
};
for(auto nitr : itr.feature_names)
{
auto _name = get_data_labels().at(itr.device_id).at(nitr);
counter_track::emplace(_dev_id, addendum(_name));
}
}
}
for(auto& ditr : _device_time)
{
for(auto itr = ditr.second.begin(); itr != ditr.second.end(); ++itr)
{
auto _next = std::next(itr);
if(_next == ditr.second.end()) continue;
_device_range[ditr.first].emplace(((*_next / 2) + (*itr / 2)));
}
}
}
for(auto& ditr : _device_range)
{
auto _dev_id = ditr.first;
auto _values = std::vector<rocm_feature_value>{};
auto _ts_sorted_data = _device_data[_dev_id];
std::sort(_ts_sorted_data.begin(), _ts_sorted_data.end(),
[](auto* _l, auto* _r) { return _l->exit < _r->exit; });
for(const auto& itr : ditr.second)
{
auto _v = _get_events(_ts_sorted_data, itr);
uint64_t _ts = itr;
for(auto* vitr : _v)
{
size_t _n = vitr->feature_values.size();
if(_values.empty())
{
_values.reserve(_n);
for(size_t i = 0; i < _n; ++i)
{
_values.emplace_back(vitr->feature_values.at(i));
}
}
else
{
for(size_t i = 0; i < _n; ++i)
{
#ifdef __GNUC__
# pragma GCC diagnostic push
# pragma GCC diagnostic ignored "-Wdouble-promotion"
#endif
auto _plus = [](auto& _lhs, auto&& _rhs) { _lhs += _rhs; };
std::visit(_plus, _values.at(i), vitr->feature_values.at(i));
#ifdef __GNUC__
# pragma GCC diagnostic pop
#endif
}
}
}
for(size_t i = 0; i < _values.size(); ++i)
{
auto _trace_counter = [_dev_id, i, _ts](auto&& _val) {
TRACE_COUNTER("kernel_hardware_counter",
counter_track::at(_dev_id, i), _ts, _val);
};
std::visit(_trace_counter, _values.at(i));
}
}
}
}
void
post_process_timemory()
{
static bool _once = false;
if(_once) return;
auto _data = rocm_data_t{};
auto _device_data = std::map<uint32_t, std::vector<rocm_event*>>{};
auto _device_fields = std::map<uint32_t, std::vector<std::string_view>>{};
auto _device_range = std::map<uint32_t, std::set<rocm_metric_type>>{};
for(size_t i = 0; i < ROCPROFSYS_MAX_THREADS; ++i)
{
auto& _v = component::rocm_data(i);
if(_v)
{
_data.reserve(_data.size() + _v->size());
for(auto& itr : *_v)
_data.emplace_back(itr);
}
}
if(_data.empty()) return;
_once = true;
std::sort(_data.begin(), _data.end());
for(auto& itr : _data)
{
_device_data[itr.device_id].emplace_back(&itr);
}
for(auto& itr : _device_data)
{
// sort according to when it exited
std::sort(itr.second.begin(), itr.second.end(),
[](auto* _lhs, auto* _rhs) { return _lhs->exit < _rhs->exit; });
}
using storage_type = typename rocm_data_tracker::storage_type;
using bundle_type = tim::lightweight_tuple<rocm_data_tracker>;
auto _info = rocm_metrics();
static auto _get_description = [&_info](std::string_view _v) {
for(auto& itr : _info)
{
if(itr.symbol().find(_v) == 0 || itr.short_description().find(_v) == 0)
{
return itr.long_description();
}
}
return std::string{};
};
struct local_event
{
rocm_event* parent = nullptr;
mutable std::vector<local_event> children = {};
ROCPROFSYS_DEFAULT_OBJECT(local_event)
explicit local_event(rocm_event* _v)
: parent{ _v }
{}
bool operator()(rocm_event* _v)
{
if(!parent) return false;
if(_v->device_id != parent->device_id) return false;
if(_v->entry > parent->entry && _v->exit <= parent->exit)
{
children.emplace_back(_v);
return true;
}
return false;
}
bool operator<(const local_event& _v) const
{
if(!parent && _v.parent) return true;
if(parent && !_v.parent) return false;
return *parent < *_v.parent;
}
void operator()(int64_t _index, scope::config _scope) const
{
if(!parent) return;
bundle_type _bundle{ parent->name, _scope };
_bundle.push(parent->queue_id)
.start()
.store(parent->feature_values.at(_index));
std::sort(children.begin(), children.end());
for(const auto& itr : children)
itr(_index, _scope);
_bundle.stop().pop(parent->queue_id);
}
};
struct local_storage
{
int64_t index = 0;
std::string metric_name = {};
std::string metric_description = {};
std::unique_ptr<storage_type> storage = {};
local_storage(uint32_t _devid, size_t _idx, std::string_view _name)
: index{ static_cast<int64_t>(_idx) }
, metric_name{ _name }
, metric_description{ _get_description(metric_name) }
{
auto _metric_name = std::string{ _name };
_metric_name = std::regex_replace(
_metric_name, std::regex{ "(.*)\\[([0-9]+)\\]" }, "$1_$2");
storage = std::make_unique<storage_type>(
tim::standalone_storage{}, index,
JOIN('-', "rocprof", "device", _devid, _metric_name));
}
void operator()(const local_event& _event, scope::config _scope) const
{
operation::set_storage<rocm_data_tracker>{}(storage.get());
_event(index, _scope);
}
void write() const
{
rocm_data_tracker::label() = metric_name;
rocm_data_tracker::description() = metric_description;
storage->write();
}
};
auto _local_data = std::map<uint32_t, std::vector<local_event>>{};
auto _scope = scope::get_default();
for(auto& ditr : _device_data)
{
ROCPROFSYS_VERBOSE_F(1, "Post-processing %zu entries for device %u...\n",
ditr.second.size(), ditr.first);
auto _storage = std::vector<local_storage>{};
for(auto& itr : ditr.second)
{
auto _n = itr->feature_names.size();
if(_n > _storage.size())
{
_storage.reserve(_n);
for(size_t i = _storage.size(); i < _n; ++i)
_storage.emplace_back(
ditr.first, i,
get_data_labels().at(ditr.first).at(itr->feature_names.at(i)));
}
}
auto& _local = _local_data[ditr.first];
_local.reserve(ditr.second.size());
double _avg = 0.0;
for(auto& itr : ditr.second)
{
if(_local.empty() || itr->entry >= _local.back().parent->exit)
{
_local.emplace_back(itr);
}
else
{
size_t _n = 0;
bool _found = false;
for(auto litr = _local.rbegin(); litr != _local.rend(); ++litr)
{
++_n;
if((*litr)(itr))
{
_found = true;
break;
}
}
if(!_found) _local.emplace_back(itr);
_avg += _n;
}
}
ROCPROFSYS_VERBOSE_F(3, "Average # of iterations before match: %.1f\n",
_avg / ditr.second.size() * 100.0);
for(auto& sitr : _storage)
{
for(auto& itr : _local)
sitr(itr, _scope);
}
for(auto& itr : _storage)
itr.write();
}
tim::trait::runtime_enabled<rocprofsys::rocprofiler::rocm_data_tracker>::set(false);
}
} // namespace
void
post_process()
{
if(get_use_perfetto()) post_process_perfetto();
if(get_use_timemory())
{
auto _manager = tim::manager::master_instance();
if(_manager)
{
_manager->add_cleanup("rocprofiler", &post_process_timemory);
}
else
{
post_process_timemory();
}
}
}
} // namespace rocprofiler
} // namespace rocprofsys
source/lib/rocprof-sys/library/roctracer.cpp
-967
View File
@@ -1,967 +0,0 @@
// MIT License
//
// Copyright (c) 2022-2024 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 "library/roctracer.hpp"
#include "binary/analysis.hpp"
#include "core/components/fwd.hpp"
#include "core/concepts.hpp"
#include "core/config.hpp"
#include "core/debug.hpp"
#include "core/locking.hpp"
#include "library/components/category_region.hpp"
#include "library/runtime.hpp"
#include "library/sampling.hpp"
#include "library/thread_data.hpp"
#include "library/thread_info.hpp"
#include "library/tracing.hpp"
#include <timemory/backends/cpu.hpp>
#include <timemory/backends/threading.hpp>
#include <timemory/hash/types.hpp>
#include <timemory/utility/types.hpp>
#include <atomic>
#include <chrono>
#include <cstdint>
#include <tuple>
#include <roctracer_ext.h>
#include <roctracer_hip.h>
#include <roctracer_roctx.h>
#if ROCPROFSYS_HIP_VERSION < 50300
# include <roctracer_hcc.h>
#endif
#define AMD_INTERNAL_BUILD 1
#include <roctracer_hsa.h>
#if __has_include(<hip/amd_detail/hip_prof_str.h>) || (defined(ROCPROFSYS_USE_HIP) && ROCPROFSYS_USE_HIP > 0)
# include <hip/amd_detail/hip_prof_str.h>
# define ROCPROFSYS_HIP_API_ARGS 1
#else
# define ROCPROFSYS_HIP_API_ARGS 0
#endif
TIMEMORY_DEFINE_API(roctracer)
namespace rocprofsys
{
namespace
{
template <typename Tp>
auto&
roctracer_type_mutex()
{
return tim::type_mutex<Tp, category::roctracer, max_supported_threads,
locking::atomic_mutex>();
}
std::string
hip_api_string(hip_api_id_t id, const hip_api_data_t* data)
{
#if ROCPROFSYS_HIP_API_ARGS > 0
std::string _v = hipApiString(id, data);
if(_v.empty()) return _v;
auto _pbeg = _v.find('(');
if(_pbeg == std::string::npos) return _v;
auto _pend = _v.find_last_of(')');
if(_pend == std::string::npos || _pbeg >= _pend) return _v;
auto _n = (_pend - _pbeg - 1);
return _v.substr(_pbeg + 1, _n);
#else
tim::consume_parameters(id, data);
#endif
}
int&
get_current_device()
{
static thread_local int _v = 1;
return _v;
}
std::unordered_set<uint64_t>&
get_roctracer_kernels()
{
static auto _v = std::unordered_set<uint64_t>{};
return _v;
}
auto&
get_roctracer_hip_data(int64_t _tid = threading::get_id())
{
using data_t = std::unordered_map<uint64_t, roctracer_hip_bundle_t>;
using thread_data_t = thread_data<data_t, category::roctracer>;
return thread_data_t::instance(construct_on_thread{ _tid });
}
std::unordered_map<uint64_t, const char*>&
get_roctracer_key_data()
{
static auto _v = std::unordered_map<uint64_t, const char*>{};
return _v;
}
std::unordered_map<uint64_t, int64_t>&
get_roctracer_tid_data()
{
static auto _v = std::unordered_map<uint64_t, int64_t>{};
return _v;
}
auto&
get_hip_activity_callbacks(int64_t _tid = threading::get_id())
{
using thread_data_t =
thread_data<std::vector<std::function<void()>>, category::roctracer>;
return thread_data_t::instance(construct_on_thread{ _tid });
}
size_t
get_hip_activity_callbacks_size()
{
using thread_data_t =
thread_data<std::vector<std::function<void()>>, category::roctracer>;
return thread_data_t::size();
}
using hip_activity_mutex_t = std::decay_t<decltype(get_hip_activity_callbacks())>;
using key_data_mutex_t = std::decay_t<decltype(get_roctracer_key_data())>;
auto&
get_hip_activity_mutex(int64_t _tid = threading::get_id())
{
return tim::type_mutex<hip_activity_mutex_t, category::roctracer,
max_supported_threads, locking::atomic_mutex>(
_tid % max_supported_threads);
}
} // namespace
//
int64_t
get_clock_skew()
{
static auto _use = tim::get_env("ROCPROFSYS_USE_ROCTRACER_CLOCK_SKEW", true);
if(!_use) return 0;
static auto _v = []() {
auto _gpu_now = []() {
uint64_t _ts = 0;
roctracer_get_timestamp(&_ts);
return _ts;
};
// discard (warm-up)
(void) tracing::get_clock_skew(_gpu_now, 1);
auto _diff = tracing::get_clock_skew(_gpu_now, 10);
ROCPROFSYS_BASIC_VERBOSE(1, "CPU/HIP timestamp skew: %li (used: %s)\n", _diff,
_use ? "yes" : "no");
return _diff;
}();
return _v;
}
// HSA API callback function
void
hsa_api_callback(uint32_t domain, uint32_t cid, const void* callback_data, void* arg)
{
if(get_state() != State::Active || !trait::runtime_enabled<comp::roctracer>::get())
return;
ROCPROFSYS_SCOPED_THREAD_STATE(ThreadState::Internal);
(void) arg;
const hsa_api_data_t* data = reinterpret_cast<const hsa_api_data_t*>(callback_data);
ROCPROFSYS_CONDITIONAL_PRINT_F(
get_debug() && get_verbose() >= 2, "<%-30s id(%u)\tcorrelation_id(%lu) %s>\n",
roctracer_op_string(domain, cid, 0), cid, data->correlation_id,
(data->phase == ACTIVITY_API_PHASE_ENTER) ? "on-enter" : "on-exit");
static thread_local int64_t begin_timestamp = 0;
switch(cid)
{
case HSA_API_ID_hsa_init:
case HSA_API_ID_hsa_shut_down:
case HSA_API_ID_hsa_agent_get_exception_policies:
case HSA_API_ID_hsa_agent_get_info:
case HSA_API_ID_hsa_amd_agent_iterate_memory_pools:
case HSA_API_ID_hsa_amd_agent_memory_pool_get_info:
case HSA_API_ID_hsa_amd_coherency_get_type:
case HSA_API_ID_hsa_amd_memory_pool_get_info:
case HSA_API_ID_hsa_amd_pointer_info:
case HSA_API_ID_hsa_amd_pointer_info_set_userdata:
case HSA_API_ID_hsa_amd_profiling_async_copy_enable:
case HSA_API_ID_hsa_amd_profiling_get_async_copy_time:
case HSA_API_ID_hsa_amd_profiling_get_dispatch_time:
case HSA_API_ID_hsa_amd_profiling_set_profiler_enabled:
case HSA_API_ID_hsa_cache_get_info:
case HSA_API_ID_hsa_code_object_get_info:
case HSA_API_ID_hsa_code_object_get_symbol:
case HSA_API_ID_hsa_code_object_get_symbol_from_name:
case HSA_API_ID_hsa_code_object_reader_create_from_memory:
case HSA_API_ID_hsa_code_symbol_get_info:
case HSA_API_ID_hsa_executable_create_alt:
case HSA_API_ID_hsa_executable_freeze:
case HSA_API_ID_hsa_executable_get_info:
case HSA_API_ID_hsa_executable_get_symbol:
case HSA_API_ID_hsa_executable_get_symbol_by_name:
case HSA_API_ID_hsa_executable_symbol_get_info:
case HSA_API_ID_hsa_extension_get_name:
case HSA_API_ID_hsa_ext_image_data_get_info:
case HSA_API_ID_hsa_ext_image_data_get_info_with_layout:
case HSA_API_ID_hsa_ext_image_get_capability:
case HSA_API_ID_hsa_ext_image_get_capability_with_layout:
case HSA_API_ID_hsa_isa_get_exception_policies:
case HSA_API_ID_hsa_isa_get_info:
case HSA_API_ID_hsa_isa_get_info_alt:
case HSA_API_ID_hsa_isa_get_round_method:
case HSA_API_ID_hsa_region_get_info:
case HSA_API_ID_hsa_system_extension_supported:
case HSA_API_ID_hsa_system_get_extension_table:
case HSA_API_ID_hsa_system_get_info:
case HSA_API_ID_hsa_system_get_major_extension_table:
case HSA_API_ID_hsa_wavefront_get_info: break;
default:
{
if(data->phase == ACTIVITY_API_PHASE_ENTER)
{
begin_timestamp = comp::wall_clock::record();
}
else
{
const auto* _name = roctracer_op_string(domain, cid, 0);
const auto end_timestamp = (cid == HSA_API_ID_hsa_shut_down)
? begin_timestamp
: comp::wall_clock::record();
if(begin_timestamp > end_timestamp) return;
if(get_use_perfetto())
{
uint64_t _beg_ts = begin_timestamp;
uint64_t _end_ts = end_timestamp;
tracing::push_perfetto_ts(category::rocm_hsa{}, _name, _beg_ts,
[&](::perfetto::EventContext ctx) {
if(config::get_perfetto_annotations())
{
tracing::add_perfetto_annotation(
ctx, "begin_ns", _beg_ts);
}
});
tracing::pop_perfetto_ts(category::rocm_hsa{}, _name, _end_ts,
[&](::perfetto::EventContext ctx) {
if(config::get_perfetto_annotations())
{
tracing::add_perfetto_annotation(
ctx, "end_ns", _end_ts);
}
});
}
if(get_use_timemory())
{
auto _beg_ns = begin_timestamp;
auto _end_ns = end_timestamp;
if(tasking::roctracer::get_task_group().pool())
tasking::roctracer::get_task_group().exec(
[_name, _beg_ns, _end_ns]() {
roctracer_hsa_bundle_t _bundle{ _name };
_bundle.start()
.store(std::plus<double>{},
static_cast<double>(_end_ns - _beg_ns))
.stop();
});
}
// timemory is disabled in this callback because collecting data in this
// thread causes strange segmentation faults
}
}
}
}
void
hsa_activity_callback(uint32_t op, const void* vrecord, void* arg)
{
const auto* record = static_cast<const activity_record_t*>(vrecord);
if(get_state() != State::Active || !trait::runtime_enabled<comp::roctracer>::get())
return;
ROCPROFSYS_SCOPED_THREAD_STATE(ThreadState::Internal);
auto&& _protect = comp::roctracer::protect_flush_activity();
(void) _protect;
static const char* copy_op_name = "hsa_async_copy";
static const char* dispatch_op_name = "hsa_dispatch";
static const char* barrier_op_name = "hsa_barrier";
const char** _name = nullptr;
switch(op)
{
case HSA_OP_ID_DISPATCH: _name = &dispatch_op_name; break;
case HSA_OP_ID_COPY: _name = &copy_op_name; break;
case HSA_OP_ID_BARRIER: _name = &barrier_op_name; break;
default: break;
}
ROCPROFSYS_CI_FAIL(_name == nullptr, "Error! HSA operation type not handled: %u\n",
op);
if(!_name) return;
auto _beg_ns = record->begin_ns + get_clock_skew();
auto _end_ns = record->end_ns + get_clock_skew();
if(get_use_perfetto())
{
uint64_t _beg = _beg_ns;
uint64_t _end = _end_ns;
tracing::push_perfetto_ts(
category::device_hsa{}, *_name, _beg, [&](::perfetto::EventContext ctx) {
if(config::get_perfetto_annotations())
{
tracing::add_perfetto_annotation(ctx, "begin_ns", _beg);
}
});
tracing::pop_perfetto_ts(
category::device_hsa{}, *_name, _end, [&](::perfetto::EventContext ctx) {
if(config::get_perfetto_annotations())
{
tracing::add_perfetto_annotation(ctx, "end_ns", _end);
}
});
}
auto _func = [_beg_ns, _end_ns, _name]() {
if(get_use_timemory())
{
roctracer_hsa_bundle_t _bundle{ *_name };
_bundle.start()
.store(std::plus<double>{}, static_cast<double>(_end_ns - _beg_ns))
.stop();
}
};
if(tasking::roctracer::get_task_group().pool())
tasking::roctracer::get_task_group().exec(_func);
// timemory is disabled in this callback because collecting data in this thread
// causes strange segmentation faults
tim::consume_parameters(arg);
}
void
hip_exec_activity_callbacks(int64_t _tid)
{
// guard against initialization of structure when trying to exec
if(static_cast<size_t>(_tid) >= get_hip_activity_callbacks_size()) return;
// ROCPROFSYS_ROCTRACER_CALL(roctracer_flush_activity());
locking::atomic_lock _lk{ get_hip_activity_mutex(_tid) };
auto& _async_ops = get_hip_activity_callbacks(_tid);
if(!_async_ops) return;
for(auto& itr : *_async_ops)
{
if(itr) itr();
}
_async_ops->clear();
}
namespace
{
thread_local std::unordered_map<size_t, size_t> gpu_crit_cids = {};
}
void
roctx_api_callback(uint32_t domain, uint32_t cid, const void* callback_data,
void* /*arg*/)
{
if(get_state() != State::Active || !trait::runtime_enabled<comp::roctracer>::get())
return;
ROCPROFSYS_SCOPED_THREAD_STATE(ThreadState::Internal);
if(domain != ACTIVITY_DOMAIN_ROCTX) return;
static auto _range_map = std::unordered_map<roctx_range_id_t, std::string>{};
static auto _range_lock = locking::atomic_mutex{};
const auto* _data = reinterpret_cast<const roctx_api_data_t*>(callback_data);
static thread_local auto _range_stack = std::vector<std::string>{};
switch(cid)
{
case ROCTX_API_ID_roctxRangePushA:
{
if(_data->args.message)
{
auto& itr = _range_stack.emplace_back(std::string{ _data->args.message });
component::category_region<category::rocm_roctx>::start(itr.c_str());
}
break;
}
case ROCTX_API_ID_roctxRangePop:
{
if(!_range_stack.empty())
{
auto& itr = _range_stack.back();
component::category_region<category::rocm_roctx>::stop(itr.c_str());
_range_stack.pop_back();
}
else
{
ROCPROFSYS_THROW("Error! roctxRangePop stack is empty! Expected "
"roctxRangePush/roctxRangePop on same thread\n");
}
break;
}
case ROCTX_API_ID_roctxRangeStartA:
{
{
locking::atomic_lock _lk{ _range_lock, std::defer_lock };
if(!_lk.owns_lock()) _lk.lock();
_range_map.emplace(roctx_range_id_t{ _data->args.id },
std::string{ _data->args.message });
}
component::category_region<category::rocm_roctx>::start(_data->args.message);
break;
}
case ROCTX_API_ID_roctxRangeStop:
{
std::string_view _message = {};
{
locking::atomic_lock _lk{ _range_lock, std::defer_lock };
if(!_lk.owns_lock()) _lk.lock();
auto itr = _range_map.find(roctx_range_id_t{ _data->args.id });
ROCPROFSYS_CI_THROW(itr == _range_map.end(),
"Error! could not find range with id %lu\n",
_data->args.id);
if(itr == _range_map.end())
{
ROCPROFSYS_VERBOSE(0, "Warning! could not find range with id %lu\n",
_data->args.id);
return;
}
else
{
_message = itr->second;
}
}
if(!_message.empty())
{
component::category_region<category::rocm_roctx>::stop(_message.data());
}
break;
}
case ROCTX_API_ID_roctxMarkA:
{
if(_data->args.message)
{
component::category_region<category::rocm_roctx>::mark(
_data->args.message);
}
break;
}
default: break;
}
}
// HIP API callback function
void
hip_api_callback(uint32_t domain, uint32_t cid, const void* callback_data, void* arg)
{
if(get_state() != State::Active || !trait::runtime_enabled<comp::roctracer>::get())
return;
ROCPROFSYS_SCOPED_THREAD_STATE(ThreadState::Internal);
assert(domain == ACTIVITY_DOMAIN_HIP_API);
const char* op_name = roctracer_op_string(domain, cid, 0);
if(op_name == nullptr) op_name = hip_api_name(cid);
if(op_name == nullptr) return;
assert(std::string{ op_name } == std::string{ hip_api_name(cid) });
switch(cid)
{
case HIP_API_ID___hipPushCallConfiguration:
case HIP_API_ID___hipPopCallConfiguration:
case HIP_API_ID_hipDeviceEnablePeerAccess:
#if ROCPROFSYS_HIP_VERSION_MAJOR > 4 || \
(ROCPROFSYS_HIP_VERSION_MAJOR == 4 && ROCPROFSYS_HIP_VERSION_MINOR >= 3)
case HIP_API_ID_hipImportExternalMemory:
case HIP_API_ID_hipDestroyExternalMemory:
#endif
return;
default: break;
}
const hip_api_data_t* data = reinterpret_cast<const hip_api_data_t*>(callback_data);
ROCPROFSYS_CONDITIONAL_PRINT_F(
get_debug() && get_verbose() >= 2, "<%-30s id(%u)\tcorrelation_id(%lu) %s>\n",
op_name, cid, data->correlation_id,
(data->phase == ACTIVITY_API_PHASE_ENTER) ? "on-enter" : "on-exit");
int64_t _ts = comp::wall_clock::record();
auto _tid = threading::get_id();
uint64_t _crit_cid = 0;
uint64_t _parent_crit_cid = 0;
uint32_t _depth = 0;
auto _roct_cid = data->correlation_id;
auto& _device_id = get_current_device();
if(data->phase == ACTIVITY_API_PHASE_ENTER)
{
if(cid == HIP_API_ID_hipSetDevice)
get_current_device() =
reinterpret_cast<int>(data->args.hipSetDevice.deviceId) + 1;
const char* _name = nullptr;
switch(cid)
{
case HIP_API_ID_hipLaunchKernel:
{
_name = hipKernelNameRefByPtr(data->args.hipLaunchKernel.function_address,
data->args.hipLaunchKernel.stream);
break;
}
case HIP_API_ID_hipLaunchCooperativeKernel:
{
_name =
hipKernelNameRefByPtr(data->args.hipLaunchCooperativeKernel.f,
data->args.hipLaunchCooperativeKernel.stream);
if(!_name)
{
_name =
hipKernelNameRefByPtr(data->args.hipLaunchKernel.function_address,
data->args.hipLaunchKernel.stream);
}
break;
}
case HIP_API_ID_hipHccModuleLaunchKernel:
{
_name = hipKernelNameRef(data->args.hipHccModuleLaunchKernel.f);
break;
}
case HIP_API_ID_hipModuleLaunchKernel:
{
_name = hipKernelNameRef(data->args.hipModuleLaunchKernel.f);
break;
}
case HIP_API_ID_hipExtModuleLaunchKernel:
{
_name = hipKernelNameRef(data->args.hipExtModuleLaunchKernel.f);
break;
}
case HIP_API_ID_hipExtLaunchKernel:
{
_name =
hipKernelNameRefByPtr(data->args.hipExtLaunchKernel.function_address,
data->args.hipLaunchKernel.stream);
break;
}
default: break;
}
if(_name != nullptr)
{
if(get_use_perfetto() || get_use_timemory() || get_use_rocm_smi())
{
locking::atomic_lock _lk{ roctracer_type_mutex<key_data_mutex_t>() };
get_roctracer_key_data().emplace(_roct_cid, _name);
get_roctracer_tid_data().emplace(_roct_cid, _tid);
}
}
std::tie(_crit_cid, _parent_crit_cid, _depth) = create_cpu_cid_entry();
if(get_use_perfetto())
{
static auto _compact_annotations =
config::get_setting_value<bool>(
"ROCPROFSYS_PERFETTO_COMPACT_ROCTRACER_ANNOTATIONS")
.value_or(false);
static auto _enable_backtraces =
config::get_setting_value<bool>("ROCPROFSYS_ROCTRACER_HIP_API_BACKTRACE")
.value_or(false);
constexpr size_t bt_stack_depth = 16;
constexpr size_t bt_ignore_depth = 3;
constexpr bool bt_with_signal_frame = true;
using backtrace_entry_vec_t = std::vector<tim::unwind::processed_entry>;
auto _bt_data = std::optional<backtrace_entry_vec_t>{};
if(_enable_backtraces && config::get_perfetto_annotations())
{
auto _backtrace = tim::get_unw_stack<bt_stack_depth, bt_ignore_depth,
bt_with_signal_frame>();
_bt_data = backtrace_entry_vec_t{};
_bt_data->reserve(_backtrace.size());
for(auto itr : _backtrace)
{
if(itr)
{
if(auto _val = binary::lookup_ipaddr_entry<false>(itr->address());
_val)
{
_bt_data->emplace_back(std::move(*_val));
}
}
}
}
auto _api_id = static_cast<hip_api_id_t>(cid);
tracing::push_perfetto_ts(
category::rocm_hip{}, op_name, _ts,
::perfetto::Flow::ProcessScoped(_roct_cid),
[&](::perfetto::EventContext ctx) {
if(config::get_perfetto_annotations())
{
tracing::add_perfetto_annotation(ctx, "begin_ns", _ts);
tracing::add_perfetto_annotation(ctx, "cid", _crit_cid);
tracing::add_perfetto_annotation(ctx, "pcid", _parent_crit_cid);
tracing::add_perfetto_annotation(ctx, "device", _device_id);
tracing::add_perfetto_annotation(ctx, "tid", _tid);
tracing::add_perfetto_annotation(ctx, "depth", _depth);
tracing::add_perfetto_annotation(ctx, "corr_id", _roct_cid);
if(_compact_annotations)
{
tracing::add_perfetto_annotation(
ctx, "args", hip_api_string(_api_id, data));
}
else
{
auto _args = std::string{ hip_api_string(_api_id, data) };
if(!_args.empty())
{
for(auto itr : tim::delimit(_args, ","))
{
if(itr.empty()) continue;
auto _bpos = itr.find_first_not_of(' ');
auto _epos = itr.find_last_not_of(' ');
if(_epos > _bpos)
itr = itr.substr(_bpos, (_epos - _bpos) + 1);
auto _pos = itr.find('=');
if(_pos != std::string::npos)
tracing::add_perfetto_annotation(
ctx, itr.substr(0, _pos),
itr.substr(_pos + 1));
}
}
}
if(_enable_backtraces && _bt_data && !_bt_data->empty())
{
const std::string _unk = "??";
size_t _bt_cnt = 0;
for(const auto& itr : *_bt_data)
{
const auto* _func =
(itr.name.empty()) ? &_unk : &itr.name;
const auto* _loc =
(itr.location.empty()) ? &_unk : &itr.location;
auto _line = (itr.lineno == 0) ? std::string{ "?" }
: join("", itr.lineno);
auto _entry = join("", demangle(*_func), " @ ",
join(':', *_loc, _line));
if(_bt_cnt < 10)
{
// Prepend zero for better ordering in UI.
// Only one zero is ever necessary since stack depth
// is limited to 16.
tracing::add_perfetto_annotation(
ctx, join("", "frame#0", _bt_cnt++), _entry);
}
else
{
tracing::add_perfetto_annotation(
ctx, join("", "frame#", _bt_cnt++), _entry);
}
}
}
}
});
}
if(get_use_timemory())
{
auto itr = get_roctracer_hip_data()->emplace(
_roct_cid, roctracer_hip_bundle_t{ op_name });
if(itr.second)
{
itr.first->second.start();
}
else if(itr.first != get_roctracer_hip_data()->end())
{
itr.first->second.stop();
get_roctracer_hip_data()->erase(itr.first);
}
}
hip_exec_activity_callbacks(_tid);
}
else if(data->phase == ACTIVITY_API_PHASE_EXIT)
{
hip_exec_activity_callbacks(_tid);
if(get_use_perfetto())
{
tracing::pop_perfetto_ts(
category::rocm_hip{}, op_name, _ts, [&](::perfetto::EventContext ctx) {
if(config::get_perfetto_annotations())
{
tracing::add_perfetto_annotation(ctx, "end_ns", _ts);
}
});
}
if(get_use_timemory())
{
auto _stop = [&_roct_cid](int64_t _tid_v) {
auto& _data = get_roctracer_hip_data(_tid_v);
auto itr = _data->find(_roct_cid);
if(itr != get_roctracer_hip_data()->end())
{
itr->second.stop();
_data->erase(itr);
return true;
}
return false;
};
if(!_stop(_tid))
{
for(size_t i = 0; i < thread_info::get_peak_num_threads(); ++i)
{
if(_stop(i)) break;
}
}
}
}
tim::consume_parameters(arg);
}
// Activity tracing callback
void
hip_activity_callback(const char* begin, const char* end, void* arg)
{
if(get_state() != State::Active || !trait::runtime_enabled<comp::roctracer>::get())
return;
ROCPROFSYS_SCOPED_THREAD_STATE(ThreadState::Internal);
auto&& _protect = comp::roctracer::protect_flush_activity();
(void) _protect;
if(!trait::runtime_enabled<comp::roctracer>::get()) return;
static auto _kernel_names = std::unordered_map<const char*, std::string>{};
static auto _indexes = std::unordered_map<uint64_t, int>{};
static auto _skip_barrier_packets =
config::get_setting_value<bool>("ROCPROFSYS_ROCTRACER_DISCARD_BARRIERS")
.value_or(false);
const roctracer_record_t* record = reinterpret_cast<const roctracer_record_t*>(begin);
const roctracer_record_t* end_record =
reinterpret_cast<const roctracer_record_t*>(end);
auto&& _advance_record = [&record]() {
ROCPROFSYS_ROCTRACER_CALL(roctracer_next_record(record, &record));
};
while(record < end_record)
{
// make sure every iteration advances regardless of where return point happens
scope::destructor _next_dtor{ _advance_record };
// ROCPROFSYS_CI will enable these asserts and should fail if something relevant
// changes
assert(HIP_OP_ID_DISPATCH == 0);
assert(HIP_OP_ID_COPY == 1);
assert(HIP_OP_ID_BARRIER == 2);
if(record->domain == ACTIVITY_DOMAIN_HSA_OPS)
{
hsa_activity_callback(record->op, record, arg);
continue;
}
if(record->domain != ACTIVITY_DOMAIN_HIP_OPS) continue;
if(record->op > HIP_OP_ID_BARRIER) continue;
if(_skip_barrier_packets && record->op == HIP_OP_ID_BARRIER) continue;
const char* op_name =
roctracer_op_string(record->domain, record->op, record->kind);
auto _ns_skew = get_clock_skew();
uint64_t _beg_ns = record->begin_ns + _ns_skew;
uint64_t _end_ns = record->end_ns + _ns_skew;
auto _roct_cid = record->correlation_id;
auto& _keys = get_roctracer_key_data();
auto& _tids = get_roctracer_tid_data();
int64_t _tid = 0; // thread id
int32_t _devid = record->device_id; // device id
int64_t _queid = record->queue_id; // queue id
uintptr_t _queue = 0; // Host queue (stream)
const char* _name = nullptr;
bool _found = false;
{
locking::atomic_lock _lk{ roctracer_type_mutex<key_data_mutex_t>() };
if(_tids.find(_roct_cid) != _tids.end())
{
_found = true;
_tid = _tids.at(_roct_cid);
auto itr = _keys.find(_roct_cid);
if(itr != _keys.end()) _name = itr->second;
}
}
if(_name == nullptr && op_name == nullptr) continue;
if(_name == nullptr) _name = op_name;
static auto _op_id_names =
std::array<const char*, 3>{ "DISPATCH", "COPY", "BARRIER" };
if(_end_ns < _beg_ns)
{
auto _verbose = []() { return get_verbose() >= 0 || get_debug(); };
static size_t _n = 0;
static size_t _nmax =
get_env<size_t>("ROCPROFSYS_ROCTRACER_DISCARD_INVALID", 0);
if(_nmax == 0) std::swap(_end_ns, _beg_ns);
ROCPROFSYS_WARNING_IF_F(
_n < _nmax && _verbose(),
"%4zu :: Discarding kernel roctracer activity record which ended before "
"it started :: %-20s :: %-20s :: cid=%lu, time_ns=(%12lu:%12lu) "
"delta=%li, device=%d, queue=%lu, pid=%u, tid=%lu, op=%s\n",
_n, op_name, _name, record->correlation_id, _beg_ns, _end_ns,
(static_cast<int64_t>(_end_ns) - static_cast<int64_t>(_beg_ns)), _devid,
_queid, record->process_id, _tid, _op_id_names.at(record->op));
ROCPROFSYS_WARNING_IF_F(
_nmax > 0 && _n == _nmax && _verbose(),
"Suppressing future messages about discarding kernel roctracer activity "
"record which ended before it started. Set "
"ROCPROFSYS_ROCTRACER_DISCARD_INVALID=N to increase/decrease the number "
"of messages. If N is set to 0, data will be included after swapping the "
"begin and end values\n");
if(_end_ns < _beg_ns)
{
++_n;
continue;
}
}
// execute this on this thread bc of how perfetto visualization works
if(get_use_perfetto())
{
if(_kernel_names.find(_name) == _kernel_names.end())
_kernel_names.emplace(_name, tim::demangle(_name));
auto _track_desc = [](int32_t _device_id, int64_t _queue_id) {
if(config::get_perfetto_roctracer_per_stream())
return JOIN("", "HIP Activity Device ", _device_id, ", Queue ",
_queue_id);
return JOIN("", "HIP Activity Device ", _device_id);
};
const auto _track = tracing::get_perfetto_track(
category::device_hip{}, _track_desc, _devid,
(get_perfetto_roctracer_per_stream()) ? _queid : 0);
assert(_end_ns >= _beg_ns);
tracing::push_perfetto_track(
category::device_hip{}, _kernel_names.at(_name).c_str(), _track, _beg_ns,
::perfetto::Flow::ProcessScoped(_roct_cid),
[&](::perfetto::EventContext ctx) {
if(config::get_perfetto_annotations())
{
tracing::add_perfetto_annotation(ctx, "begin_ns", _beg_ns);
tracing::add_perfetto_annotation(ctx, "end_ns", _end_ns);
tracing::add_perfetto_annotation(ctx, "corr_id", _roct_cid);
tracing::add_perfetto_annotation(ctx, "device", _devid);
tracing::add_perfetto_annotation(ctx, "queue", _queid);
tracing::add_perfetto_annotation(ctx, "tid", _tid);
tracing::add_perfetto_annotation(
ctx, "stream", JOIN("", "0x", std::hex, _queue));
tracing::add_perfetto_annotation(ctx, "op",
_op_id_names.at(record->op));
}
});
tracing::pop_perfetto_track(category::device_hip{}, "", _track, _end_ns);
}
if(_found && _name != nullptr && get_use_timemory())
{
auto _func = [_beg_ns, _end_ns, _name]() {
roctracer_hip_bundle_t _bundle{ _name };
_bundle.start()
.store(std::plus<double>{}, static_cast<double>(_end_ns - _beg_ns))
.stop()
.get<comp::wall_clock>([&](comp::wall_clock* wc) {
wc->set_value(_end_ns - _beg_ns);
wc->set_accum(_end_ns - _beg_ns);
return wc;
});
_bundle.pop();
};
auto& _async_ops = get_hip_activity_callbacks(_tid);
locking::atomic_lock _lk{ get_hip_activity_mutex(_tid) };
_async_ops->emplace_back(std::move(_func));
}
}
// ensures that all the updates are written
if(get_use_perfetto()) ::perfetto::TrackEvent::Flush();
}
bool&
roctracer_is_init()
{
static bool _v = tim::get_env("ROCPROFSYS_ROCTRACER_IS_INIT", false);
return _v;
}
bool&
roctracer_is_setup()
{
static bool _v = false;
return _v;
}
using roctracer_functions_t = std::vector<std::pair<std::string, std::function<void()>>>;
roctracer_functions_t&
roctracer_setup_routines()
{
static auto _v = roctracer_functions_t{};
return _v;
}
roctracer_functions_t&
roctracer_shutdown_routines()
{
static auto _v = roctracer_functions_t{};
return _v;
}
} // namespace rocprofsys
source/lib/rocprof-sys/library/roctracer.hpp
-89
View File
@@ -1,89 +0,0 @@
// MIT License
//
// Copyright (c) 2022-2024 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 "core/config.hpp"
#include "core/debug.hpp"
#include "core/hip_runtime.hpp"
#include "core/perfetto.hpp"
#include "library/components/roctracer.hpp"
#include "library/ptl.hpp"
#include <iostream>
#include <memory>
// Macro to check ROC-tracer calls status
#define ROCPROFSYS_ROCTRACER_CALL(call) \
{ \
ROCPROFSYS_DEBUG_F(#call); \
int err = call; \
if(err != 0) \
{ \
ROCPROFSYS_PRINT_F("%s in: %s\n", roctracer_error_string(), #call); \
} \
}
namespace rocprofsys
{
using roctracer_hip_bundle_t =
tim::component_bundle<category::rocm_hip, comp::roctracer_data, comp::wall_clock>;
using roctracer_hsa_bundle_t =
tim::component_bundle<category::rocm_hsa, comp::roctracer_data>;
using roctracer_functions_t = std::vector<std::pair<std::string, std::function<void()>>>;
// HSA API callback function
void
hsa_api_callback(uint32_t domain, uint32_t cid, const void* callback_data, void* arg);
void
hsa_activity_callback(uint32_t op, const void* record, void* arg);
void
hip_exec_activity_callbacks(int64_t _tid);
// HIP API callback function
void
hip_api_callback(uint32_t domain, uint32_t cid, const void* callback_data, void* arg);
void
roctx_api_callback(uint32_t domain, uint32_t cid, const void* callback_data, void* arg);
// Activity tracing callback
void
hip_activity_callback(const char* begin, const char* end, void*);
bool&
roctracer_is_init();
bool&
roctracer_is_setup();
int64_t
get_clock_skew();
roctracer_functions_t&
roctracer_setup_routines();
roctracer_functions_t&
roctracer_shutdown_routines();
} // namespace rocprofsys
source/lib/rocprof-sys/library/runtime.hpp
-1
View File
@@ -33,7 +33,6 @@
#include "library/components/mpi_gotcha.hpp"
#include "library/components/numa_gotcha.hpp"
#include "library/components/pthread_gotcha.hpp"
#include "library/components/roctracer.hpp"
#include "library/thread_data.hpp"
#include <timemory/backends/threading.hpp>