AI/rocm-systems
2
0
Forka 0
Codice Problemi Pull Request Actions Pacchetti Progetti Rilasci Wiki Attività
Files
1e612a5e52103fd134b48a092dee3086d22c2f16
rocm-systems/source/lib/rocprofiler-sdk/hsa/async_copy.cpp
T
Benjamin Welton 1e612a5e52 Wait for all memory copies to complete before allowing destruction (#725) 
* Wait for all mem copies to complete before destroying.

* Update source/lib/rocprofiler-sdk/hsa/async_copy.cpp

Co-authored-by: Ammar ELWazir <ammar.elwazir@amd.com>

* Update async_copy.cpp

---------

Co-authored-by: Ammar ELWazir <ammar.elwazir@amd.com>
2024-04-02 08:22:37 -05:00

691 righe
25 KiB
C++
Originale Blame Cronologia

// MIT License
//
// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to deal
// in the Software without restriction, including without limitation the rights
// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
// copies of the Software, and to permit persons to whom the Software is
// furnished to do so, subject to the following conditions:
//
// The above copyright notice and this permission notice shall be included in
// all copies or substantial portions of the Software.
//
// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
// THE SOFTWARE.
#include "lib/rocprofiler-sdk/hsa/async_copy.hpp"
#include "lib/common/defines.hpp"
#include "lib/common/static_object.hpp"
#include "lib/common/utility.hpp"
#include "lib/rocprofiler-sdk/agent.hpp"
#include "lib/rocprofiler-sdk/buffer.hpp"
#include "lib/rocprofiler-sdk/context/context.hpp"
#include "lib/rocprofiler-sdk/hsa/details/ostream.hpp"
#include "lib/rocprofiler-sdk/hsa/hsa.hpp"
#include "lib/rocprofiler-sdk/hsa/utils.hpp"
#include <rocprofiler-sdk/fwd.h>
#include <rocprofiler-sdk/hsa/api_id.h>
#include <rocprofiler-sdk/hsa/table_id.h>
#include <glog/logging.h>
#include <hsa/amd_hsa_signal.h>
#include <hsa/hsa.h>
#include <cstdlib>
#include <type_traits>
#define ROCP_HSA_TABLE_CALL(SEVERITY, EXPR) \
auto ROCPROFILER_VARIABLE(rocp_hsa_table_call_, __LINE__) = (EXPR); \
LOG_IF(SEVERITY, ROCPROFILER_VARIABLE(rocp_hsa_table_call_, __LINE__) != HSA_STATUS_SUCCESS) \
<< #EXPR << " returned non-zero status code " \
<< ROCPROFILER_VARIABLE(rocp_hsa_table_call_, __LINE__) << " "
#if defined(ROCPROFILER_CI)
# define ROCP_CI_LOG_IF(NON_CI_LEVEL, ...) LOG_IF(FATAL, __VA_ARGS__)
# define ROCP_CI_LOG(NON_CI_LEVEL, ...) LOG(FATAL)
#else
# define ROCP_CI_LOG_IF(NON_CI_LEVEL, ...) LOG_IF(NON_CI_LEVEL, __VA_ARGS__)
# define ROCP_CI_LOG(NON_CI_LEVEL, ...) LOG(NON_CI_LEVEL)
#endif
#define ROCPROFILER_LIB_ROCPROFILER_HSA_ASYNC_COPY_CPP_IMPL 1
// template specializations
#include "hsa.def.cpp"
namespace rocprofiler
{
namespace hsa
{
namespace async_copy
{
namespace
{
using context_t = context::context;
using context_array_t = common::container::small_vector<const context_t*>;
using external_corr_id_map_t = std::unordered_map<const context_t*, rocprofiler_user_data_t>;
template <size_t OpIdx>
struct async_copy_info;
#define SPECIALIZE_ASYNC_COPY_INFO(DIRECTION) \
template <> \
struct async_copy_info<ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_##DIRECTION> \
{ \
static constexpr auto operation_idx = ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_##DIRECTION; \
static constexpr auto name = #DIRECTION; \
};
SPECIALIZE_ASYNC_COPY_INFO(NONE)
SPECIALIZE_ASYNC_COPY_INFO(HOST_TO_HOST)
SPECIALIZE_ASYNC_COPY_INFO(HOST_TO_DEVICE)
SPECIALIZE_ASYNC_COPY_INFO(DEVICE_TO_HOST)
SPECIALIZE_ASYNC_COPY_INFO(DEVICE_TO_DEVICE)
#undef SPECIALIZE_ASYNC_COPY_INFO
template <size_t Idx, size_t... IdxTail>
const char*
name_by_id(const uint32_t id, std::index_sequence<Idx, IdxTail...>)
{
if(Idx == id) return async_copy_info<Idx>::name;
if constexpr(sizeof...(IdxTail) > 0)
return name_by_id(id, std::index_sequence<IdxTail...>{});
else
return nullptr;
}
template <size_t Idx, size_t... IdxTail>
uint32_t
id_by_name(const char* name, std::index_sequence<Idx, IdxTail...>)
{
if(std::string_view{async_copy_info<Idx>::name} == std::string_view{name})
return async_copy_info<Idx>::operation_idx;
if constexpr(sizeof...(IdxTail) > 0)
return id_by_name(name, std::index_sequence<IdxTail...>{});
else
return ROCPROFILER_HSA_AMD_EXT_API_ID_NONE;
}
template <size_t... Idx>
void
get_ids(std::vector<uint32_t>& _id_list, std::index_sequence<Idx...>)
{
auto _emplace = [](auto& _vec, uint32_t _v) {
if(_v < static_cast<uint32_t>(ROCPROFILER_HSA_AMD_EXT_API_ID_LAST)) _vec.emplace_back(_v);
};
(_emplace(_id_list, async_copy_info<Idx>::operation_idx), ...);
}
template <size_t... Idx>
void
get_names(std::vector<const char*>& _name_list, std::index_sequence<Idx...>)
{
auto _emplace = [](auto& _vec, const char* _v) {
if(_v != nullptr && strnlen(_v, 1) > 0) _vec.emplace_back(_v);
};
(_emplace(_name_list, async_copy_info<Idx>::name), ...);
}
bool
context_filter(const context::context* ctx)
{
return (ctx->buffered_tracer &&
(ctx->buffered_tracer->domains(ROCPROFILER_BUFFER_TRACING_MEMORY_COPY)));
}
constexpr auto null_rocp_agent_id =
rocprofiler_agent_id_t{.handle = std::numeric_limits<uint64_t>::max()};
struct async_copy_data
{
hsa_signal_t orig_signal = {};
hsa_signal_t rocp_signal = {};
rocprofiler_thread_id_t tid = common::get_tid();
rocprofiler_agent_id_t dst_agent = null_rocp_agent_id;
rocprofiler_agent_id_t src_agent = null_rocp_agent_id;
rocprofiler_memory_copy_operation_t direction = ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_NONE;
context::correlation_id* correlation_id = nullptr;
context::context_array_t contexts = {};
external_corr_id_map_t extern_corr_ids = {};
};
struct active_signals
{
static hsa_signal_t get_signal()
{
static hsa_signal_t signal = []() {
hsa_signal_t _signal;
ROCP_HSA_TABLE_CALL(ERROR,
get_core_table()->hsa_signal_create_fn(0, 0, nullptr, &_signal));
return _signal;
}();
return signal;
}
void sync() const
{
if(_is_set)
{
ROCP_HSA_TABLE_CALL(
ERROR,
get_core_table()->hsa_signal_wait_relaxed_fn(
get_signal(), HSA_SIGNAL_CONDITION_EQ, 0, -1, HSA_WAIT_STATE_ACTIVE));
}
}
void fetch_sub(int v)
{
_is_set = true;
get_core_table()->hsa_signal_subtract_relaxed_fn(get_signal(), v);
}
void fetch_add(int v)
{
_is_set = true;
get_core_table()->hsa_signal_add_relaxed_fn(get_signal(), v);
}
~active_signals()
{
if(_is_set)
{
sync();
ROCP_HSA_TABLE_CALL(ERROR, get_core_table()->hsa_signal_destroy_fn(get_signal()));
}
}
std::atomic<bool> _is_set{false};
};
active_signals*
get_active_signals()
{
static auto* _v = common::static_object<active_signals>::construct();
return _v;
}
template <typename Tp, typename Up>
constexpr Tp*
convert_hsa_handle(Up _hsa_object)
{
static_assert(!std::is_pointer<Up>::value, "pass opaque struct");
static_assert(!std::is_pointer<Tp>::value, "pass non-pointer type");
// NOLINTNEXTLINE(performance-no-int-to-ptr)
return reinterpret_cast<Tp*>(_hsa_object.handle);
}
bool
async_copy_handler(hsa_signal_value_t signal_value, void* arg)
{
static auto sysclock_period = []() -> uint64_t {
constexpr auto nanosec = 1000000000UL;
uint64_t sysclock_hz = 0;
ROCP_HSA_TABLE_CALL(ERROR,
get_core_table()->hsa_system_get_info_fn(
HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY, &sysclock_hz));
return (nanosec / sysclock_hz);
}();
auto* _data = static_cast<async_copy_data*>(arg);
auto copy_time = hsa_amd_profiling_async_copy_time_t{};
auto copy_time_status = get_amd_ext_table()->hsa_amd_profiling_get_async_copy_time_fn(
_data->rocp_signal, &copy_time);
// normalize
copy_time.start *= sysclock_period;
copy_time.end *= sysclock_period;
// if we encounter this in CI, it will cause test to fail
ROCP_CI_LOG_IF(ERROR, copy_time_status == HSA_STATUS_SUCCESS && copy_time.end < copy_time.start)
<< "hsa_amd_profiling_get_async_copy_time for returned async times where the end time ("
<< copy_time.end << ") was less than the start time (" << copy_time.start << ")";
// get the contexts that were active when the signal was created
const auto& ctxs = _data->contexts;
// we need to decrement this reference count at the end of the functions
auto* _corr_id = _data->correlation_id;
if(copy_time_status == HSA_STATUS_SUCCESS && !ctxs.empty())
{
const auto& _extern_corr_ids = _data->extern_corr_ids;
for(const auto* itr : ctxs)
{
auto* _buffer = buffer::get_buffer(
itr->buffered_tracer->buffer_data.at(ROCPROFILER_BUFFER_TRACING_MEMORY_COPY));
// go ahead and create the correlation id value since we expect at least one of these
// domains will require it
auto _corr_id_v =
rocprofiler_correlation_id_t{.internal = 0, .external = context::null_user_data};
if(_corr_id)
{
_corr_id_v.internal = _corr_id->internal;
_corr_id_v.external = _extern_corr_ids.at(itr);
}
if(itr->buffered_tracer->domains(ROCPROFILER_BUFFER_TRACING_MEMORY_COPY))
{
if(copy_time_status == HSA_STATUS_SUCCESS)
{
auto record = rocprofiler_buffer_tracing_memory_copy_record_t{
sizeof(rocprofiler_buffer_tracing_memory_copy_record_t),
ROCPROFILER_BUFFER_TRACING_MEMORY_COPY,
_data->direction,
_corr_id_v,
copy_time.start * sysclock_period,
copy_time.end * sysclock_period,
_data->dst_agent,
_data->src_agent};
CHECK_NOTNULL(_buffer)->emplace(ROCPROFILER_BUFFER_CATEGORY_TRACING,
ROCPROFILER_BUFFER_TRACING_MEMORY_COPY,
record);
}
}
}
}
auto* orig_amd_signal = convert_hsa_handle<amd_signal_t>(_data->orig_signal);
// Original intercepted signal completion
if(orig_amd_signal)
{
// NOLINTNEXTLINE(performance-no-int-to-ptr)
auto* rocp_amd_signal = convert_hsa_handle<amd_signal_t>(_data->rocp_signal);
std::tie(orig_amd_signal->start_ts, orig_amd_signal->end_ts) =
std::tie(rocp_amd_signal->start_ts, rocp_amd_signal->end_ts);
const hsa_signal_value_t new_value =
get_core_table()->hsa_signal_load_relaxed_fn(_data->orig_signal) - 1;
LOG_IF(ERROR, signal_value != new_value) << "bad original signal value in " << __FUNCTION__;
get_core_table()->hsa_signal_store_screlease_fn(_data->orig_signal, signal_value);
}
if(signal_value == 0)
{
ROCP_HSA_TABLE_CALL(ERROR, get_core_table()->hsa_signal_destroy_fn(_data->rocp_signal));
delete _data;
}
get_active_signals()->fetch_sub(1);
if(_corr_id) _corr_id->sub_ref_count();
return (signal_value > 0);
}
enum async_copy_id
{
async_copy_id = ROCPROFILER_HSA_AMD_EXT_API_ID_hsa_amd_memory_async_copy,
async_copy_on_engine_id = ROCPROFILER_HSA_AMD_EXT_API_ID_hsa_amd_memory_async_copy_on_engine,
async_copy_rect_id = ROCPROFILER_HSA_AMD_EXT_API_ID_hsa_amd_memory_async_copy_rect,
};
template <size_t TableIdx, size_t OpIdx>
auto&
get_next_dispatch()
{
using function_t = typename hsa_api_meta<TableIdx, OpIdx>::function_type;
static function_t _v = nullptr;
return _v;
}
template <size_t Idx>
struct arg_indices;
#define HSA_ASYNC_COPY_DEFINE_ARG_INDICES( \
ENUM_ID, DST_AGENT_IDX, SRC_AGENT_IDX, COMPLETION_SIGNAL_IDX) \
template <> \
struct arg_indices<ENUM_ID> \
{ \
static constexpr auto dst_agent_idx = DST_AGENT_IDX; \
static constexpr auto src_agent_idx = SRC_AGENT_IDX; \
static constexpr auto completion_signal_idx = COMPLETION_SIGNAL_IDX; \
};
HSA_ASYNC_COPY_DEFINE_ARG_INDICES(async_copy_id, 1, 3, 7)
HSA_ASYNC_COPY_DEFINE_ARG_INDICES(async_copy_on_engine_id, 1, 3, 7)
HSA_ASYNC_COPY_DEFINE_ARG_INDICES(async_copy_rect_id, 5, 5, 9)
template <typename FuncT, typename ArgsT, size_t... Idx>
decltype(auto)
invoke(FuncT&& _func, ArgsT&& _args, std::index_sequence<Idx...>)
{
return std::forward<FuncT>(_func)(std::get<Idx>(_args)...);
}
template <size_t TableIdx, size_t OpIdx, typename... Args>
hsa_status_t
async_copy_impl(Args... args)
{
using meta_type = hsa_api_meta<TableIdx, OpIdx>;
constexpr auto N = sizeof...(Args);
auto&& _tied_args = std::tie(args...);
auto ctxs = context::get_active_contexts(context_filter);
// no active contexts so just execute original
if(ctxs.empty())
{
return invoke(get_next_dispatch<TableIdx, OpIdx>(),
std::move(_tied_args),
std::make_index_sequence<N>{});
}
// determine the direction of the memory copy
auto _direction = ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_NONE;
auto _src_agent_id = rocprofiler_agent_id_t{};
auto _dst_agent_id = rocprofiler_agent_id_t{};
{
// indices in the tuple with references to the arguments
constexpr auto dst_agent_idx = arg_indices<OpIdx>::dst_agent_idx;
constexpr auto src_agent_idx = arg_indices<OpIdx>::src_agent_idx;
// extract the completion signal argument and the destination hsa_agent_t
auto _hsa_dst_agent = std::get<dst_agent_idx>(_tied_args);
auto _hsa_src_agent = std::get<src_agent_idx>(_tied_args);
// map the hsa agents to rocprofiler agents
auto _rocp_dst_agent = agent::get_rocprofiler_agent(_hsa_dst_agent);
auto _rocp_src_agent = agent::get_rocprofiler_agent(_hsa_src_agent);
if(_rocp_dst_agent && _rocp_src_agent)
{
_src_agent_id = _rocp_src_agent->id;
_dst_agent_id = _rocp_dst_agent->id;
if(_rocp_src_agent->type == ROCPROFILER_AGENT_TYPE_CPU)
{
if(_rocp_dst_agent->type == ROCPROFILER_AGENT_TYPE_CPU)
_direction = ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_HOST_TO_HOST;
else if(_rocp_dst_agent->type == ROCPROFILER_AGENT_TYPE_GPU)
_direction = ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_HOST_TO_DEVICE;
else
{
ROCP_CI_LOG(WARNING)
<< meta_type::name
<< " had an unhandled destination type: " << _rocp_dst_agent->type;
}
}
else if(_rocp_src_agent->type == ROCPROFILER_AGENT_TYPE_GPU)
{
if(_rocp_dst_agent->type == ROCPROFILER_AGENT_TYPE_CPU)
_direction = ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_DEVICE_TO_HOST;
else if(_rocp_dst_agent->type == ROCPROFILER_AGENT_TYPE_GPU)
_direction = ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_DEVICE_TO_DEVICE;
else
{
ROCP_CI_LOG(WARNING)
<< meta_type::name
<< " had an unhandled destination type: " << _rocp_dst_agent->type;
}
}
else
{
ROCP_CI_LOG(WARNING) << meta_type::name
<< " had an unhandled source type: " << _rocp_dst_agent->type;
}
}
else
{
LOG_IF(ERROR, !_rocp_src_agent)
<< "failed to find source rocprofiler agent for hsa agent with handle="
<< _hsa_src_agent.handle;
LOG_IF(ERROR, !_rocp_dst_agent)
<< "failed to find destination rocprofiler agent for hsa agent with handle="
<< _hsa_dst_agent.handle;
}
}
// remove any contexts which do not wish to trace this memory copy direction
ctxs.erase(std::remove_if(ctxs.begin(),
ctxs.end(),
[_direction](const context_t* ctx) {
return !ctx->buffered_tracer->domains(
ROCPROFILER_BUFFER_TRACING_MEMORY_COPY, _direction);
}),
ctxs.end());
// if no contexts remain, execute as usual
if(ctxs.empty())
{
return invoke(get_next_dispatch<TableIdx, OpIdx>(),
std::move(_tied_args),
std::make_index_sequence<N>{});
}
// at this point, we want to install our own signal handler
auto* _data = new async_copy_data{};
_data->tid = common::get_tid();
_data->dst_agent = _dst_agent_id;
_data->src_agent = _src_agent_id;
_data->direction = _direction;
_data->contexts = ctxs; // avoid using move in case code below accidentally uses ctxs
constexpr auto completion_signal_idx = arg_indices<OpIdx>::completion_signal_idx;
auto& _completion_signal = std::get<completion_signal_idx>(_tied_args);
const hsa_signal_value_t _completion_signal_val =
get_core_table()->hsa_signal_load_scacquire_fn(_completion_signal);
{
const uint32_t num_consumers = 0;
const hsa_agent_t* consumers = nullptr;
auto _status = get_core_table()->hsa_signal_create_fn(
_completion_signal_val, num_consumers, consumers, &_data->rocp_signal);
if(_status != HSA_STATUS_SUCCESS)
{
LOG(ERROR) << "hsa_signal_create returned non-zero error code " << _status;
delete _data;
return invoke(get_next_dispatch<TableIdx, OpIdx>(),
std::move(_tied_args),
std::make_index_sequence<N>{});
}
else
{
get_active_signals()->fetch_add(1);
}
}
{
auto _status = get_amd_ext_table()->hsa_amd_signal_async_handler_fn(_data->rocp_signal,
HSA_SIGNAL_CONDITION_LT,
_completion_signal_val,
async_copy_handler,
_data);
if(_status != HSA_STATUS_SUCCESS)
{
LOG(ERROR) << "hsa_amd_signal_async_handler returned non-zero error code " << _status;
ROCP_HSA_TABLE_CALL(ERROR, get_core_table()->hsa_signal_destroy_fn(_data->rocp_signal))
<< ":: failed to destroy signal after async handler failed";
get_active_signals()->fetch_sub(1);
delete _data;
return invoke(get_next_dispatch<TableIdx, OpIdx>(),
std::move(_tied_args),
std::make_index_sequence<N>{});
}
}
_data->correlation_id = context::get_latest_correlation_id();
auto& extern_corr_ids = _data->extern_corr_ids;
// increase the reference count to denote that this correlation id is being used in a kernel
if(_data->correlation_id)
{
extern_corr_ids.reserve(_data->contexts.size()); // reserve for performance
for(const auto* ctx : _data->contexts)
extern_corr_ids.emplace(ctx,
ctx->correlation_tracer.external_correlator.get(_data->tid));
_data->correlation_id->add_ref_count();
}
_data->orig_signal = _completion_signal;
_completion_signal = _data->rocp_signal;
return invoke(
get_next_dispatch<TableIdx, OpIdx>(), std::move(_tied_args), std::make_index_sequence<N>{});
}
template <size_t TableIdx, size_t OpIdx, typename RetT, typename... Args>
auto get_async_copy_impl(RetT (*)(Args...))
{
return &async_copy_impl<TableIdx, OpIdx, Args...>;
}
template <size_t TableIdx, size_t OpIdx>
void
async_copy_save(hsa_amd_ext_table_t* _orig, uint64_t _tbl_instance)
{
static_assert(
std::is_same<hsa_amd_ext_table_t, typename hsa_table_lookup<TableIdx>::type>::value,
"unexpected type");
auto _meta = hsa_api_meta<TableIdx, OpIdx>{};
// original table and function
auto& _orig_table = _meta.get_table(_orig);
auto& _orig_func = _meta.get_table_func(_orig_table);
// table with copy function
auto& _copy_func = get_next_dispatch<TableIdx, OpIdx>();
LOG_IF(FATAL, _copy_func && _tbl_instance == 0)
<< _meta.name << " has non-null function pointer " << _copy_func
<< " despite this being the first instance of the library being copies";
if(!_copy_func)
{
LOG(INFO) << "copying table entry for " << _meta.name;
_copy_func = _orig_func;
}
else
{
LOG(INFO) << "skipping copying table entry for " << _meta.name << " from table instance "
<< _tbl_instance;
}
}
template <size_t TableIdx, size_t... OpIdx>
void
async_copy_save(hsa_amd_ext_table_t* _orig, uint64_t _tbl_instance, std::index_sequence<OpIdx...>)
{
static_assert(
std::is_same<hsa_amd_ext_table_t, typename hsa_table_lookup<TableIdx>::type>::value,
"unexpected type");
(async_copy_save<TableIdx, OpIdx>(_orig, _tbl_instance), ...);
}
template <size_t TableIdx, size_t OpIdx>
void
async_copy_wrap(hsa_amd_ext_table_t* _orig)
{
static_assert(
std::is_same<hsa_amd_ext_table_t, typename hsa_table_lookup<TableIdx>::type>::value,
"unexpected type");
auto _meta = hsa_api_meta<TableIdx, OpIdx>{};
auto& _table = _meta.get_table(_orig);
auto& _func = _meta.get_table_func(_table);
auto& _dispatch = get_next_dispatch<TableIdx, OpIdx>();
CHECK_NOTNULL(_dispatch);
_func = get_async_copy_impl<TableIdx, OpIdx>(_func);
}
template <size_t TableIdx, size_t... OpIdx>
void
async_copy_wrap(hsa_amd_ext_table_t* _orig, std::index_sequence<OpIdx...>)
{
static_assert(
std::is_same<hsa_amd_ext_table_t, typename hsa_table_lookup<TableIdx>::type>::value,
"unexpected type");
(async_copy_wrap<TableIdx, OpIdx>(_orig), ...);
}
using async_copy_index_seq_t =
std::index_sequence<async_copy_id, async_copy_on_engine_id, async_copy_rect_id>;
} // namespace
// check out the assembly here... this compiles to a switch statement
const char*
name_by_id(uint32_t id)
{
return name_by_id(id, std::make_index_sequence<ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_LAST>{});
}
uint32_t
id_by_name(const char* name)
{
return id_by_name(name,
std::make_index_sequence<ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_LAST>{});
}
std::vector<uint32_t>
get_ids()
{
auto _data = std::vector<uint32_t>{};
_data.reserve(ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_LAST);
get_ids(_data, std::make_index_sequence<ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_LAST>{});
return _data;
}
std::vector<const char*>
get_names()
{
auto _data = std::vector<const char*>{};
_data.reserve(ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_LAST);
get_names(_data, std::make_index_sequence<ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_LAST>{});
return _data;
}
} // namespace async_copy
void
async_copy_init(hsa_api_table_t* _orig, uint64_t _tbl_instance)
{
if(_orig && _orig->amd_ext_)
{
async_copy::async_copy_save<ROCPROFILER_HSA_TABLE_ID_AmdExt>(
_orig->amd_ext_, _tbl_instance, async_copy::async_copy_index_seq_t{});
auto ctxs = context::get_registered_contexts(async_copy::context_filter);
if(!ctxs.empty())
{
_orig->amd_ext_->hsa_amd_profiling_async_copy_enable_fn(true);
async_copy::async_copy_wrap<ROCPROFILER_HSA_TABLE_ID_AmdExt>(
_orig->amd_ext_, async_copy::async_copy_index_seq_t{});
}
}
}
void
async_copy_fini()
{
if(!async_copy::get_active_signals()) return;
async_copy::get_active_signals()->sync();
}
} // namespace hsa
} // namespace rocprofiler
Fai riferimento in un nuovo problema Visualizza Git Blame Copia Permalink