From 5e0a39e5d584bb2151e1acaa0ec6575a559fb869 Mon Sep 17 00:00:00 2001 From: "Jonathan R. Madsen" Date: Tue, 9 Jan 2024 11:34:46 -0600 Subject: [PATCH] Async memory copy tracing (#317) * Update samples/api_buffered_tracing/client.cpp - support ROCPROFILER_BUFFER_TRACING_MEMORY_COPY * Update include/rocprofiler-sdk/{buffer_tracing,fwd}.h - update rocprofiler_buffer_tracing_memory_copy_record_t - add ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_HOST_TO_HOST to rocprofiler_memory_copy_operation_t * Update lib/rocprofiler-sdk/context/context.* - get_registered_contexts functions (local copy) * Update tests/apps/reproducible-runtime/reproducible-runtime.cpp - include some memory allocations and memory copies for better testing * Update tests/common/serialization.hpp - update serialization save function for rocprofiler_buffer_tracing_memory_copy_record_t * Update lib/rocprofiler-sdk/hsa/hsa.* - remove stale set_callback / activity_functor_t code - forward decl hsa_api_meta - template struct hsa_api_func for getting function return type and args * Update tests/kernel-tracing/validate.py - enforce memory_copies data size - test timestamps in memory copies data - improve internal and external correlation id validation * Update lib/rocprofiler-sdk/hsa/defines.hpp - HSA_API_META_DEFINITION macro * Update lib/rocprofiler/hsa/rocprofiler-sdk/hsa/hsa.def.cpp - HSA_API_META_DEFINITION specializations for async copy functions * Add lib/rocprofiler-sdk/hsa/async_copy.{hpp,cpp} - implements buffer memory tracing * Update lib/rocprofiler-sdk/registration.cpp - invoke rocprofiler::hsa::async_copy_init * Update lib/rocprofiler-sdk/hsa/async_copy.cpp - logging improvements - improve hsa <-> rocp agent mapping * Update lib/rocprofiler-sdk/hsa/async_copy.cpp - load original signal in async signal handler before store_screlease * Update lib/rocprofiler-sdk/hsa/async_copy.cpp - use store_relaxed instead of store_screlease * Update lib/rocprofiler-sdk/hsa/async_copy.cpp - logging * Update lib/rocprofiler-sdk/hsa/async_copy.cpp - logging * Update lib/rocprofiler-sdk/hsa/async_copy.cpp - misc changes * Update lib/rocprofiler-sdk/hsa/async_copy.cpp - misc changes * Update lib/rocprofiler-sdk/hsa/async_copy.cpp - misc changes * Update lib/rocprofiler-sdk/hsa/async_copy.cpp - return function pointer instead of lambda * Update reproducible-runtime.cpp - device sync * Update tests/apps/reproducible-runtime/reproducible-runtime.cpp - use *Async variants of hipMalloc and hipMemcpy * Update lib/rocprofiler-sdk/hsa/async_copy.cpp - populate async data properly * Update tests/kernel-tracing/validate.py - verification of async copy direction * Update tests/apps/reproducible-runtime/reproducible-runtime.cpp - temporarily disable async memcpy functions * Create tests/tools - directory containing tool libraries used for collecting data in integration tests * Update tests/kernel-tracing - remove kernel-tracing-test-tool library (now rocprofiler-sdk-json-tool) - update cmake, validate.py, conftest.py accordingly * Add tests/async-copy-tracing - integration test validating async copy tracing in transpose example * Update tests/CMakeLists.txt - updates for restructuring * Revert tests/apps/reproducible-runtime - restore code to semi-original state (no memory copying) * Update tests/async-copy-tracing/validate.py - fix comment in test_async_copy_direction * Fix building tests against installation [ROCm/rocprofiler-sdk commit: 936816f762c7a44e971d95690b1011842b67a1cf] --- .../samples/api_buffered_tracing/client.cpp | 24 + .../include/rocprofiler-sdk/buffer_tracing.h | 18 +- .../source/include/rocprofiler-sdk/fwd.h | 7 +- .../lib/rocprofiler-sdk/context/context.hpp | 6 + .../lib/rocprofiler-sdk/hsa/CMakeLists.txt | 8 +- .../lib/rocprofiler-sdk/hsa/async_copy.cpp | 485 ++++++++++++++++++ .../lib/rocprofiler-sdk/hsa/async_copy.hpp | 36 ++ .../lib/rocprofiler-sdk/hsa/defines.hpp | 42 ++ .../lib/rocprofiler-sdk/hsa/hsa.def.cpp | 24 +- .../source/lib/rocprofiler-sdk/hsa/hsa.hpp | 18 + .../lib/rocprofiler-sdk/registration.cpp | 2 + projects/rocprofiler-sdk/tests/CMakeLists.txt | 6 +- .../reproducible-runtime.cpp | 1 + .../tests/async-copy-tracing/CMakeLists.txt | 47 ++ .../tests/async-copy-tracing/conftest.py | 20 + .../tests/async-copy-tracing/pytest.ini | 4 + .../tests/async-copy-tracing/validate.py | 167 ++++++ .../tests/common/serialization.hpp | 6 +- .../tests/kernel-tracing/CMakeLists.txt | 27 +- .../tests/kernel-tracing/conftest.py | 2 +- .../tests/kernel-tracing/validate.py | 91 +++- .../tests/tools/CMakeLists.txt | 37 ++ .../json-tool.cpp} | 16 +- 23 files changed, 1015 insertions(+), 79 deletions(-) create mode 100644 projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/async_copy.cpp create mode 100644 projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/async_copy.hpp create mode 100644 projects/rocprofiler-sdk/tests/async-copy-tracing/CMakeLists.txt create mode 100644 projects/rocprofiler-sdk/tests/async-copy-tracing/conftest.py create mode 100644 projects/rocprofiler-sdk/tests/async-copy-tracing/pytest.ini create mode 100644 projects/rocprofiler-sdk/tests/async-copy-tracing/validate.py create mode 100644 projects/rocprofiler-sdk/tests/tools/CMakeLists.txt rename projects/rocprofiler-sdk/tests/{kernel-tracing/kernel-tracing.cpp => tools/json-tool.cpp} (98%) diff --git a/projects/rocprofiler-sdk/samples/api_buffered_tracing/client.cpp b/projects/rocprofiler-sdk/samples/api_buffered_tracing/client.cpp index 4514918928..596626fdf2 100644 --- a/projects/rocprofiler-sdk/samples/api_buffered_tracing/client.cpp +++ b/projects/rocprofiler-sdk/samples/api_buffered_tracing/client.cpp @@ -34,6 +34,7 @@ #include "client.hpp" #include +#include #include #include #include @@ -309,6 +310,29 @@ tool_tracing_callback(rocprofiler_context_id_t context, static_cast(user_data)->emplace_back( source_location{__FUNCTION__, __FILE__, __LINE__, info.str()}); } + else if(header->category == ROCPROFILER_BUFFER_CATEGORY_TRACING && + header->kind == ROCPROFILER_BUFFER_TRACING_MEMORY_COPY) + { + auto* record = + static_cast(header->payload); + + auto info = std::stringstream{}; + + info << "src_agent_id=" << record->src_agent_id.handle + << ", dst_agent_id=" << record->dst_agent_id.handle + << ", direction=" << record->operation << ", context=" << context.handle + << ", buffer_id=" << buffer_id.handle + << ", cid=" << record->correlation_id.internal + << ", extern_cid=" << record->correlation_id.external.value + << ", kind=" << record->kind << ", start=" << record->start_timestamp + << ", stop=" << record->end_timestamp; + + if(record->start_timestamp > record->end_timestamp) + throw std::runtime_error("memory copy: start > end"); + + static_cast(user_data)->emplace_back( + source_location{__FUNCTION__, __FILE__, __LINE__, info.str()}); + } else { throw std::runtime_error{"unexpected rocprofiler_record_header_t category + kind"}; diff --git a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/buffer_tracing.h b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/buffer_tracing.h index ccaaf3f9ce..ab91a49adf 100644 --- a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/buffer_tracing.h +++ b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/buffer_tracing.h @@ -74,7 +74,7 @@ typedef struct rocprofiler_tracing_operation_t operation; ///< ::rocprofiler_marker_api_id_t rocprofiler_timestamp_t timestamp; ///< time in nanoseconds rocprofiler_thread_id_t thread_id; ///< id for thread generating this record - uint64_t marker_id; // rocprofiler_marker_id_t + uint64_t marker_id; ///< rocprofiler_marker_id_t // const char* message; // (Need Review?) } rocprofiler_buffer_tracing_marker_record_t; @@ -83,14 +83,14 @@ typedef struct */ typedef struct { - uint64_t size; ///< size of this struct - rocprofiler_buffer_tracing_kind_t kind; ///< ::ROCPROFILER_BUFFER_TRACING_MEMORY_COPY - rocprofiler_correlation_id_t correlation_id; ///< correlation ids for record - rocprofiler_timestamp_t start_timestamp; ///< start time in nanoseconds - rocprofiler_timestamp_t end_timestamp; ///< end time in nanoseconds - rocprofiler_agent_id_t agent_id; ///< agent identifier - rocprofiler_queue_id_t queue_id; ///< queue identifier - rocprofiler_kernel_id_t kernel_id; ///< kernel identifier + uint64_t size; ///< size of this struct + rocprofiler_buffer_tracing_kind_t kind; ///< ::ROCPROFILER_BUFFER_TRACING_MEMORY_COPY + rocprofiler_correlation_id_t correlation_id; ///< correlation ids for record + rocprofiler_memory_copy_operation_t operation; ///< memory copy direction + rocprofiler_timestamp_t start_timestamp; ///< start time in nanoseconds + rocprofiler_timestamp_t end_timestamp; ///< end time in nanoseconds + rocprofiler_agent_id_t dst_agent_id; ///< destination agent of copy + rocprofiler_agent_id_t src_agent_id; ///< source agent of copy } rocprofiler_buffer_tracing_memory_copy_record_t; /** diff --git a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/fwd.h b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/fwd.h index 7d66c1c0b5..99ac53aaf6 100644 --- a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/fwd.h +++ b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/fwd.h @@ -170,12 +170,13 @@ typedef enum // NOLINT(performance-enum-size) */ typedef enum // NOLINT(performance-enum-size) { - ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_NONE = 0, ///< Unknown memory copy direction - ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_DEVICE_TO_HOST, ///< Memory copy from device to host + ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_NONE = -1, ///< Unknown memory copy direction + ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_HOST_TO_HOST, ///< Memory copy from host to host ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_HOST_TO_DEVICE, ///< Memory copy from host to device + ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_DEVICE_TO_HOST, ///< Memory copy from device to host ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_DEVICE_TO_DEVICE, ///< Memory copy from device to device ROCPROFILER_BUFFER_TRACING_MEMORY_COPY_LAST, -} rocprofiler_buffer_tracing_memory_copy_operation_t; +} rocprofiler_memory_copy_operation_t; /** * @brief PC Sampling Method. diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/context/context.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/context/context.hpp index b8ad1293bb..699fd4d14a 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/context/context.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/context/context.hpp @@ -193,6 +193,12 @@ get_registered_contexts(context_array_t& data, context_filter_t filter = default context_array_t get_registered_contexts(context_filter_t filter = default_context_filter); +context_array_t& +get_registered_contexts(context_array_t& data, context_filter_t filter); + +context_array_t +get_registered_contexts(context_filter_t filter); + context_array_t& get_active_contexts(context_array_t& data, context_filter_t filter = default_context_filter); diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/CMakeLists.txt b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/CMakeLists.txt index 8e3368efa0..5e0e562260 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/CMakeLists.txt +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/CMakeLists.txt @@ -1,8 +1,11 @@ -set(ROCPROFILER_LIB_HSA_SOURCES agent_cache.cpp aql_packet.cpp code_object.cpp hsa.cpp - queue_controller.cpp queue.cpp) +# +set(ROCPROFILER_LIB_HSA_SOURCES agent_cache.cpp aql_packet.cpp async_copy.cpp + code_object.cpp hsa.cpp queue_controller.cpp queue.cpp) + set(ROCPROFILER_LIB_HSA_HEADERS agent_cache.hpp aql_packet.hpp + async_copy.hpp code_object.hpp defines.hpp hsa.hpp @@ -10,6 +13,7 @@ set(ROCPROFILER_LIB_HSA_HEADERS queue.hpp types.hpp utils.hpp) + target_sources(rocprofiler-object-library PRIVATE ${ROCPROFILER_LIB_HSA_SOURCES} ${ROCPROFILER_LIB_HSA_HEADERS}) diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/async_copy.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/async_copy.cpp new file mode 100644 index 0000000000..212f965213 --- /dev/null +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/async_copy.cpp @@ -0,0 +1,485 @@ +// 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/scope_destructor.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/utils.hpp" +#include "rocprofiler-sdk/fwd.h" +#include "rocprofiler-sdk/hsa/api_id.h" + +#include +#include +#include + +#include + +#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 +{ +using context_t = context::context; +using context_array_t = common::container::small_vector; +using external_corr_id_map_t = std::unordered_map; + +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::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 = {}; +}; + +template +constexpr Tp* +convert_hsa_handle(Up _hsa_object) +{ + static_assert(!std::is_pointer::value, "pass opaque struct"); + static_assert(!std::is_pointer::value, "pass non-pointer type"); + // NOLINTNEXTLINE(performance-no-int-to-ptr) + return reinterpret_cast(_hsa_object.handle); +} + +bool +async_copy_handler(hsa_signal_value_t signal_value, void* arg) +{ + LOG(ERROR) << "[" << __FUNCTION__ << "] invoked with signal value " << signal_value; + + static auto sysclock_period = []() -> uint64_t { + constexpr auto nanosec = 1000000000UL; + uint64_t sysclock_hz = 0; + ROCP_HSA_TABLE_CALL(ERROR, + get_table().core_->hsa_system_get_info_fn( + HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCY, &sysclock_hz)); + return (nanosec / sysclock_hz); + }(); + + auto* _data = static_cast(arg); + auto copy_time = hsa_amd_profiling_async_copy_time_t{}; + auto copy_time_status = get_table().amd_ext_->hsa_amd_profiling_get_async_copy_time_fn( + _data->rocp_signal, ©_time); + + // normalize + copy_time.start *= sysclock_period; + copy_time.end *= sysclock_period; + + LOG(ERROR) << "[" << __FUNCTION__ << "] start=" << copy_time.start << ", end=" << copy_time.end + << ", delta=" << (copy_time.end - copy_time.start) << ", period=" << 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, + _corr_id_v, + _data->direction, + 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); + } + } + } + } + + if(_corr_id) _corr_id->ref_count.fetch_sub(1); + + auto* orig_amd_signal = convert_hsa_handle(_data->orig_signal); + + // Original intercepted signal completion + if(orig_amd_signal) + { + // NOLINTNEXTLINE(performance-no-int-to-ptr) + auto* rocp_amd_signal = convert_hsa_handle(_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_table().core_->hsa_signal_load_relaxed_fn(_data->orig_signal) - 1; + + LOG_IF(ERROR, signal_value != new_value) << "bad original signal value in " << __FUNCTION__; + + get_table().core_->hsa_signal_store_screlease_fn(_data->orig_signal, signal_value); + } + + if(signal_value == 0) + { + ROCP_HSA_TABLE_CALL(ERROR, get_table().core_->hsa_signal_destroy_fn(_data->rocp_signal)); + delete _data; + } + + LOG(ERROR) << "[" << __FUNCTION__ << "] completed (signal value " << signal_value << ")"; + + return (signal_value > 0); +} + +enum async_copy_id +{ + async_copy_id = ROCPROFILER_HSA_API_ID_hsa_amd_memory_async_copy, + async_copy_on_engine_id = ROCPROFILER_HSA_API_ID_hsa_amd_memory_async_copy_on_engine, + async_copy_rect_id = ROCPROFILER_HSA_API_ID_hsa_amd_memory_async_copy_rect, +}; + +template +auto& +get_next_dispatch() +{ + using function_t = typename hsa_api_meta::function_type; + static function_t _v = nullptr; + return _v; +} + +template +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 \ + { \ + 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 +decltype(auto) +invoke(FuncT&& _func, ArgsT&& _args, std::index_sequence) +{ + return std::forward(_func)(std::get(_args)...); +} + +template +hsa_status_t +async_copy_impl(Args... args) +{ + using meta_type = hsa_api_meta; + + constexpr auto N = sizeof...(Args); + + LOG(ERROR) << "[" << __FUNCTION__ << "] started..."; + + 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(), std::move(_tied_args), std::make_index_sequence{}); + } + + // 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::dst_agent_idx; + constexpr auto src_agent_idx = arg_indices::src_agent_idx; + + // extract the completion signal argument and the destination hsa_agent_t + auto _hsa_dst_agent = std::get(_tied_args); + auto _hsa_src_agent = std::get(_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(), std::move(_tied_args), std::make_index_sequence{}); + } + + // 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::completion_signal_idx; + auto& _completion_signal = std::get(_tied_args); + const hsa_signal_value_t _completion_signal_val = + get_table().core_->hsa_signal_load_scacquire_fn(_completion_signal); + + { + const uint32_t num_consumers = 0; + const hsa_agent_t* consumers = nullptr; + auto _status = get_table().core_->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(), std::move(_tied_args), std::make_index_sequence{}); + } + } + + { + auto _status = + get_table().amd_ext_->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_table().core_->hsa_signal_destroy_fn(_data->rocp_signal)) + << ":: failed to destroy signal after async handler failed"; + + delete _data; + return invoke( + get_next_dispatch(), std::move(_tied_args), std::make_index_sequence{}); + } + } + + _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->ref_count.fetch_add(1); + } + + _data->orig_signal = _completion_signal; + _completion_signal = _data->rocp_signal; + + auto _dtor = common::scope_destructor{[]() { LOG(ERROR) << "[async_copy_impl] ... returned"; }}; + + return invoke(get_next_dispatch(), std::move(_tied_args), std::make_index_sequence{}); +} + +template +auto get_async_copy_impl(RetT (*)(Args...)) +{ + return &async_copy_impl; +} + +template +void +async_copy_save(hsa_api_table_t* _orig) +{ + auto _meta = hsa_api_meta{}; + auto& _table = _meta.get_table(_orig); + auto& _func = _meta.get_table_func(_table); + get_next_dispatch() = _func; +} + +template +void +async_copy_save(hsa_api_table_t* _orig, std::index_sequence) +{ + (async_copy_save(_orig), ...); +} + +template +void +async_copy_wrap(hsa_api_table_t* _orig) +{ + auto _meta = hsa_api_meta{}; + auto& _table = _meta.get_table(_orig); + auto& _func = _meta.get_table_func(_table); + + CHECK_NOTNULL(get_next_dispatch()); + _func = get_async_copy_impl(_func); +} + +template +void +async_copy_wrap(hsa_api_table_t* _orig, std::index_sequence) +{ + (async_copy_wrap(_orig), ...); +} + +using async_copy_index_seq_t = + std::index_sequence; +} // namespace + +void +async_copy_init(hsa_api_table_t* _orig) +{ + if(_orig) + { + async_copy_save(_orig, async_copy_index_seq_t{}); + + auto ctxs = context::get_registered_contexts(context_filter); + if(!ctxs.empty()) + { + _orig->amd_ext_->hsa_amd_profiling_async_copy_enable_fn(true); + async_copy_wrap(_orig, async_copy_index_seq_t{}); + } + } +} +} // namespace hsa +} // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/async_copy.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/async_copy.hpp new file mode 100644 index 0000000000..a99071c351 --- /dev/null +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/async_copy.hpp @@ -0,0 +1,36 @@ +// MIT License +// +// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +// THE SOFTWARE. + +#pragma once + +#include + +#include "lib/rocprofiler-sdk/hsa/hsa.hpp" + +namespace rocprofiler +{ +namespace hsa +{ +void +async_copy_init(hsa_api_table_t* _orig); +} // namespace hsa +} // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/defines.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/defines.hpp index 728f3c5dff..5a7282ecd9 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/defines.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/defines.hpp @@ -82,6 +82,48 @@ #define GET_ADDR_MEMBER_FIELDS(VAR, ...) IMPL_DETAIL_FOR_EACH(ADDR_MEMBER_, VAR, __VA_ARGS__) #define GET_NAMED_MEMBER_FIELDS(VAR, ...) IMPL_DETAIL_FOR_EACH(NAMED_MEMBER_, VAR, __VA_ARGS__) +#define HSA_API_META_DEFINITION(HSA_TABLE, HSA_API_ID, HSA_FUNC, HSA_FUNC_PTR) \ + namespace rocprofiler \ + { \ + namespace hsa \ + { \ + template <> \ + struct hsa_api_meta \ + { \ + static constexpr auto table_idx = HSA_TABLE; \ + static constexpr auto operation_idx = HSA_API_ID; \ + static constexpr auto name = #HSA_FUNC; \ + \ + using this_type = hsa_api_meta; \ + using function_type = hsa_api_func::function_type; \ + \ + static auto& get_table() { return hsa_table_lookup{}(); } \ + \ + template \ + static auto& get_table(TableT& _v) \ + { \ + return hsa_table_lookup{}(_v); \ + } \ + \ + template \ + static auto& get_table_func(TableT& _table) \ + { \ + if constexpr(std::is_pointer::value) \ + { \ + assert(_table != nullptr && "nullptr to HSA table for " #HSA_FUNC " function"); \ + return _table->HSA_FUNC_PTR; \ + } \ + else \ + { \ + return _table.HSA_FUNC_PTR; \ + } \ + } \ + \ + static auto& get_table_func() { return get_table_func(get_table()); } \ + }; \ + } \ + } + #define HSA_API_INFO_DEFINITION_0(HSA_TABLE, HSA_API_ID, HSA_FUNC, HSA_FUNC_PTR) \ namespace rocprofiler \ { \ diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa.def.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa.def.cpp index f11aa56388..5af4e7f4eb 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa.def.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa.def.cpp @@ -20,19 +20,20 @@ // OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN // THE SOFTWARE. -#if defined(ROCPROFILER_LIB_ROCPROFILER_HSA_HSA_CPP_IMPL) && \ - ROCPROFILER_LIB_ROCPROFILER_HSA_HSA_CPP_IMPL == 1 +#include "lib/common/defines.hpp" +#include "lib/rocprofiler-sdk/hsa/defines.hpp" +#include "lib/rocprofiler-sdk/hsa/hsa.hpp" -# include +#include -# include "defines.hpp" -# include "hsa.hpp" - -// clang-format off HSA_API_TABLE_LOOKUP_DEFINITION(ROCPROFILER_HSA_API_TABLE_ID_CoreApi, core_) HSA_API_TABLE_LOOKUP_DEFINITION(ROCPROFILER_HSA_API_TABLE_ID_AmdExt, amd_ext_) HSA_API_TABLE_LOOKUP_DEFINITION(ROCPROFILER_HSA_API_TABLE_ID_ImageExt, image_ext_) +#if defined(ROCPROFILER_LIB_ROCPROFILER_HSA_HSA_CPP_IMPL) && \ + ROCPROFILER_LIB_ROCPROFILER_HSA_HSA_CPP_IMPL == 1 + +// clang-format off HSA_API_INFO_DEFINITION_0(ROCPROFILER_HSA_API_TABLE_ID_CoreApi, ROCPROFILER_HSA_API_ID_hsa_init, hsa_init, hsa_init_fn) HSA_API_INFO_DEFINITION_0(ROCPROFILER_HSA_API_TABLE_ID_CoreApi, ROCPROFILER_HSA_API_ID_hsa_shut_down, hsa_shut_down, hsa_shut_down_fn) HSA_API_INFO_DEFINITION_V(ROCPROFILER_HSA_API_TABLE_ID_CoreApi, ROCPROFILER_HSA_API_ID_hsa_system_get_info, hsa_system_get_info, hsa_system_get_info_fn, attribute, value) @@ -323,6 +324,15 @@ HSA_API_INFO_DEFINITION_V(ROCPROFILER_HSA_API_TABLE_ID_AmdExt, # endif # endif +#elif defined(ROCPROFILER_LIB_ROCPROFILER_HSA_ASYNC_COPY_CPP_IMPL) && \ + ROCPROFILER_LIB_ROCPROFILER_HSA_ASYNC_COPY_CPP_IMPL == 1 + +// clang-format off +HSA_API_META_DEFINITION(ROCPROFILER_HSA_API_TABLE_ID_AmdExt, ROCPROFILER_HSA_API_ID_hsa_amd_memory_async_copy, hsa_amd_memory_async_copy, hsa_amd_memory_async_copy_fn) +HSA_API_META_DEFINITION(ROCPROFILER_HSA_API_TABLE_ID_AmdExt, ROCPROFILER_HSA_API_ID_hsa_amd_memory_async_copy_on_engine, hsa_amd_memory_async_copy_on_engine, hsa_amd_memory_async_copy_on_engine_fn) +HSA_API_META_DEFINITION(ROCPROFILER_HSA_API_TABLE_ID_AmdExt, ROCPROFILER_HSA_API_ID_hsa_amd_memory_async_copy_rect, hsa_amd_memory_async_copy_rect, hsa_amd_memory_async_copy_rect_fn) +// clang-format on + #else # error "Do not compile this file directly. It is included by lib/rocprofiler/hsa/hsa.cpp" #endif diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa.hpp index ff005daf4f..7656002e32 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa.hpp @@ -42,6 +42,24 @@ struct hsa_table_lookup; template struct hsa_api_info; +template +struct hsa_api_meta; + +template +struct hsa_api_func; + +template +struct hsa_api_func +{ + using return_type = RetT; + using args_type = std::tuple; + using function_type = RetT (*)(Args...); +}; + +template +struct hsa_api_func : hsa_api_func +{}; + template struct hsa_api_impl { diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/registration.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/registration.cpp index 86e7c82545..ac5a869ad4 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/registration.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/registration.cpp @@ -25,6 +25,7 @@ #include "lib/common/static_object.hpp" #include "lib/rocprofiler-sdk/agent.hpp" #include "lib/rocprofiler-sdk/context/context.hpp" +#include "lib/rocprofiler-sdk/hsa/async_copy.hpp" #include "lib/rocprofiler-sdk/hsa/code_object.hpp" #include "lib/rocprofiler-sdk/hsa/hsa.hpp" #include "lib/rocprofiler-sdk/hsa/queue.hpp" @@ -665,6 +666,7 @@ rocprofiler_set_api_table(const char* name, rocprofiler::agent::construct_agent_cache(hsa_api_table); rocprofiler::hsa::queue_controller_init(hsa_api_table); rocprofiler::hsa::code_object_init(hsa_api_table); + rocprofiler::hsa::async_copy_init(hsa_api_table); // any internal modifications to the HsaApiTable need to be done before we make the // copy or else those modifications will be lost when HSA API tracing is enabled diff --git a/projects/rocprofiler-sdk/tests/CMakeLists.txt b/projects/rocprofiler-sdk/tests/CMakeLists.txt index cb5f7b2406..97bc24170f 100644 --- a/projects/rocprofiler-sdk/tests/CMakeLists.txt +++ b/projects/rocprofiler-sdk/tests/CMakeLists.txt @@ -33,8 +33,12 @@ find_package(Python3 REQUIRED) # common utilities add_subdirectory(common) +# tool libraries used for data collection during integration tests +add_subdirectory(tools) + # applications used by integration tests add_subdirectory(apps) -# tool libraries +# validation tests add_subdirectory(kernel-tracing) +add_subdirectory(async-copy-tracing) diff --git a/projects/rocprofiler-sdk/tests/apps/reproducible-runtime/reproducible-runtime.cpp b/projects/rocprofiler-sdk/tests/apps/reproducible-runtime/reproducible-runtime.cpp index b3b7c11609..87ed718e44 100644 --- a/projects/rocprofiler-sdk/tests/apps/reproducible-runtime/reproducible-runtime.cpp +++ b/projects/rocprofiler-sdk/tests/apps/reproducible-runtime/reproducible-runtime.cpp @@ -161,6 +161,7 @@ run(int rank, int tid, hipStream_t stream) HIP_API_CALL(hipEventElapsedTime(&time, start, stop)); } while(static_cast(time) < nruntime); + HIP_API_CALL(hipStreamSynchronize(stream)); HIP_API_CALL(hipEventDestroy(start)); HIP_API_CALL(hipEventDestroy(stop)); diff --git a/projects/rocprofiler-sdk/tests/async-copy-tracing/CMakeLists.txt b/projects/rocprofiler-sdk/tests/async-copy-tracing/CMakeLists.txt new file mode 100644 index 0000000000..3f6bdc18ca --- /dev/null +++ b/projects/rocprofiler-sdk/tests/async-copy-tracing/CMakeLists.txt @@ -0,0 +1,47 @@ +# +# +# +cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR) + +project( + rocprofiler-tests-async-copy-tracing + LANGUAGES CXX + VERSION 0.0.0) + +find_package(rocprofiler-sdk REQUIRED) + +if(ROCPROFILER_MEMCHECK_PRELOAD_ENV) + set(PRELOAD_ENV + "${ROCPROFILER_MEMCHECK_PRELOAD_ENV}:$") +else() + set(PRELOAD_ENV "LD_PRELOAD=$") +endif() + +add_test(NAME test-async-copy-tracing-execute COMMAND $) + +set_tests_properties( + test-async-copy-tracing-execute + PROPERTIES + TIMEOUT + 45 + LABELS + "integration-tests" + ENVIRONMENT + "${PRELOAD_ENV};HSA_TOOLS_LIB=$;ROCPROFILER_TOOL_OUTPUT_FILE=async-copy-tracing-test.json" + FAIL_REGULAR_EXPRESSION + "threw an exception") + +foreach(FILENAME validate.py pytest.ini conftest.py) + configure_file(${CMAKE_CURRENT_SOURCE_DIR}/${FILENAME} + ${CMAKE_CURRENT_BINARY_DIR}/${FILENAME} COPYONLY) +endforeach() + +add_test(NAME test-async-copy-tracing-validate + COMMAND ${Python3_EXECUTABLE} ${CMAKE_CURRENT_BINARY_DIR}/validate.py --input + ${CMAKE_CURRENT_BINARY_DIR}/async-copy-tracing-test.json) + +set_tests_properties( + test-async-copy-tracing-validate + PROPERTIES TIMEOUT 45 LABELS "integration-tests" DEPENDS + test-async-copy-tracing-execute FAIL_REGULAR_EXPRESSION + "threw an exception") diff --git a/projects/rocprofiler-sdk/tests/async-copy-tracing/conftest.py b/projects/rocprofiler-sdk/tests/async-copy-tracing/conftest.py new file mode 100644 index 0000000000..3b5b57a8f8 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/async-copy-tracing/conftest.py @@ -0,0 +1,20 @@ +#!/usr/bin/env python3 + +import json +import pytest + + +def pytest_addoption(parser): + parser.addoption( + "--input", + action="store", + default="async-copy-tracing-test.json", + help="Input JSON", + ) + + +@pytest.fixture +def input_data(request): + filename = request.config.getoption("--input") + with open(filename, "r") as inp: + return json.load(inp) diff --git a/projects/rocprofiler-sdk/tests/async-copy-tracing/pytest.ini b/projects/rocprofiler-sdk/tests/async-copy-tracing/pytest.ini new file mode 100644 index 0000000000..899bc38408 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/async-copy-tracing/pytest.ini @@ -0,0 +1,4 @@ + +[pytest] +addopts = --durations=20 -ras -vv +testpaths = validate.py diff --git a/projects/rocprofiler-sdk/tests/async-copy-tracing/validate.py b/projects/rocprofiler-sdk/tests/async-copy-tracing/validate.py new file mode 100644 index 0000000000..e0cfa316f6 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/async-copy-tracing/validate.py @@ -0,0 +1,167 @@ +#!/usr/bin/env python3 + +import sys +import pytest + + +# helper function +def node_exists(name, data, min_len=1): + assert name in data + assert data[name] is not None + assert len(data[name]) >= min_len + + +def test_data_structure(input_data): + """verify minimum amount of expected data is present""" + data = input_data + + node_exists("rocprofiler-sdk-json-tool", data) + node_exists("agents", data["rocprofiler-sdk-json-tool"]) + node_exists("call_stack", data["rocprofiler-sdk-json-tool"]) + node_exists("callback_records", data["rocprofiler-sdk-json-tool"]) + node_exists("buffer_records", data["rocprofiler-sdk-json-tool"]) + + node_exists("names", data["rocprofiler-sdk-json-tool"]["callback_records"]) + node_exists("code_objects", data["rocprofiler-sdk-json-tool"]["callback_records"]) + node_exists("kernel_symbols", data["rocprofiler-sdk-json-tool"]["callback_records"]) + node_exists("hsa_api_traces", data["rocprofiler-sdk-json-tool"]["callback_records"]) + + node_exists("names", data["rocprofiler-sdk-json-tool"]["buffer_records"]) + node_exists("kernel_dispatches", data["rocprofiler-sdk-json-tool"]["buffer_records"]) + node_exists("memory_copies", data["rocprofiler-sdk-json-tool"]["buffer_records"], 4) + node_exists("hsa_api_traces", data["rocprofiler-sdk-json-tool"]["buffer_records"]) + + +def test_timestamps(input_data): + data = input_data + + cb_start = {} + cb_end = {} + for itr in data["rocprofiler-sdk-json-tool"]["callback_records"]["hsa_api_traces"]: + cid = itr["record"]["correlation_id"]["internal"] + phase = itr["record"]["phase"] + if phase == 1: + cb_start[cid] = itr["timestamp"] + elif phase == 2: + cb_end[cid] = itr["timestamp"] + assert cb_start[cid] <= itr["timestamp"] + else: + assert phase == 1 or phase == 2 + + for itr in data["rocprofiler-sdk-json-tool"]["buffer_records"]["hsa_api_traces"]: + assert itr["start_timestamp"] <= itr["end_timestamp"] + + for itr in data["rocprofiler-sdk-json-tool"]["buffer_records"]["memory_copies"]: + assert itr["start_timestamp"] <= itr["end_timestamp"] + + for itr in data["rocprofiler-sdk-json-tool"]["buffer_records"]["kernel_dispatches"]: + assert itr["start_timestamp"] < itr["end_timestamp"] + assert itr["correlation_id"]["internal"] > 0 + assert itr["correlation_id"]["external"] > 0 + + api_start = cb_start[itr["correlation_id"]["internal"]] + api_end = cb_end[itr["correlation_id"]["internal"]] + assert api_start < itr["start_timestamp"] + assert api_end <= itr["end_timestamp"] + + +def test_internal_correlation_ids(input_data): + data = input_data + + api_corr_ids = [] + for itr in data["rocprofiler-sdk-json-tool"]["callback_records"]["hsa_api_traces"]: + api_corr_ids.append(itr["record"]["correlation_id"]["internal"]) + + for itr in data["rocprofiler-sdk-json-tool"]["buffer_records"]["hsa_api_traces"]: + api_corr_ids.append(itr["correlation_id"]["internal"]) + + api_corr_ids_sorted = sorted(api_corr_ids) + api_corr_ids_unique = list(set(api_corr_ids)) + + for itr in data["rocprofiler-sdk-json-tool"]["buffer_records"]["kernel_dispatches"]: + assert itr["correlation_id"]["internal"] in api_corr_ids_unique + + for itr in data["rocprofiler-sdk-json-tool"]["buffer_records"]["memory_copies"]: + assert itr["correlation_id"]["internal"] in api_corr_ids_unique + + len_corr_id_unq = len(api_corr_ids_unique) + assert len(api_corr_ids) != len_corr_id_unq + assert max(api_corr_ids_sorted) == len_corr_id_unq + + +def test_external_correlation_ids(input_data): + data = input_data + + extern_corr_ids = [] + for itr in data["rocprofiler-sdk-json-tool"]["callback_records"]["hsa_api_traces"]: + assert itr["record"]["correlation_id"]["external"] > 0 + assert itr["record"]["thread_id"] == itr["record"]["correlation_id"]["external"] + extern_corr_ids.append(itr["record"]["correlation_id"]["external"]) + + extern_corr_ids = list(set(sorted(extern_corr_ids))) + for itr in data["rocprofiler-sdk-json-tool"]["buffer_records"]["hsa_api_traces"]: + assert itr["correlation_id"]["external"] > 0 + assert itr["thread_id"] == itr["correlation_id"]["external"] + assert itr["thread_id"] in extern_corr_ids + assert itr["correlation_id"]["external"] in extern_corr_ids + + for itr in data["rocprofiler-sdk-json-tool"]["buffer_records"]["kernel_dispatches"]: + assert itr["correlation_id"]["external"] > 0 + assert itr["correlation_id"]["external"] in extern_corr_ids + + for itr in data["rocprofiler-sdk-json-tool"]["buffer_records"]["memory_copies"]: + assert itr["correlation_id"]["external"] > 0 + assert itr["correlation_id"]["external"] in extern_corr_ids + + +def test_kernel_ids(input_data): + data = input_data + + symbol_info = {} + for itr in data["rocprofiler-sdk-json-tool"]["callback_records"]["kernel_symbols"]: + phase = itr["record"]["phase"] + payload = itr["payload"] + kern_id = payload["kernel_id"] + + assert phase == 1 or phase == 2 + assert kern_id > 0 + if phase == 1: + assert len(payload["kernel_name"]) > 0 + symbol_info[kern_id] = payload + elif phase == 2: + assert payload["kernel_id"] in symbol_info.keys() + assert payload["kernel_name"] == symbol_info[kern_id]["kernel_name"] + + for itr in data["rocprofiler-sdk-json-tool"]["buffer_records"]["kernel_dispatches"]: + assert itr["kernel_id"] in symbol_info.keys() + + +def test_async_copy_direction(input_data): + data = input_data + + # Direction values: + # -1 == ??? (unknown) + # 0 == H2H (host to host) + # 1 == H2D (host to device) + # 2 == D2H (device to host) + # 3 == D2D (device to device) + async_dir_cnt = dict([(idx, 0) for idx in range(-1, 4)]) + for itr in data["rocprofiler-sdk-json-tool"]["buffer_records"]["memory_copies"]: + op_id = itr["operation"] + assert op_id > 0 + assert op_id < 3 + async_dir_cnt[op_id] += 1 + + # in the transpose test which generates the input file, + # two threads and the main thread (so three threads total) + # each perform one H2D + one D2H memory copy + assert async_dir_cnt[-1] == 0 + assert async_dir_cnt[0] == 0 + assert async_dir_cnt[1] == 6 + assert async_dir_cnt[2] == 6 + assert async_dir_cnt[3] == 0 + + +if __name__ == "__main__": + exit_code = pytest.main(["-x", __file__] + sys.argv[1:]) + sys.exit(exit_code) diff --git a/projects/rocprofiler-sdk/tests/common/serialization.hpp b/projects/rocprofiler-sdk/tests/common/serialization.hpp index c5a61c614d..0f0ede1cee 100644 --- a/projects/rocprofiler-sdk/tests/common/serialization.hpp +++ b/projects/rocprofiler-sdk/tests/common/serialization.hpp @@ -240,11 +240,11 @@ save(ArchiveT& ar, rocprofiler_buffer_tracing_memory_copy_record_t data) SAVE_DATA_FIELD(size); SAVE_DATA_FIELD(kind); SAVE_DATA_FIELD(correlation_id); + SAVE_DATA_FIELD(operation); SAVE_DATA_FIELD(start_timestamp); SAVE_DATA_FIELD(end_timestamp); - SAVE_DATA_FIELD(agent_id); - SAVE_DATA_FIELD(queue_id); - SAVE_DATA_FIELD(kernel_id); + SAVE_DATA_FIELD(dst_agent_id); + SAVE_DATA_FIELD(src_agent_id); } template diff --git a/projects/rocprofiler-sdk/tests/kernel-tracing/CMakeLists.txt b/projects/rocprofiler-sdk/tests/kernel-tracing/CMakeLists.txt index 6fe854790f..7483cc25ca 100644 --- a/projects/rocprofiler-sdk/tests/kernel-tracing/CMakeLists.txt +++ b/projects/rocprofiler-sdk/tests/kernel-tracing/CMakeLists.txt @@ -10,30 +10,11 @@ project( find_package(rocprofiler-sdk REQUIRED) -add_library(kernel-tracing-test-tool SHARED) -target_sources(kernel-tracing-test-tool PRIVATE kernel-tracing.cpp) -target_link_libraries( - kernel-tracing-test-tool - PRIVATE rocprofiler::rocprofiler rocprofiler::cereal rocprofiler::tests-build-flags - rocprofiler::tests-common-library) -set_target_properties( - kernel-tracing-test-tool - PROPERTIES INSTALL_RPATH "\$ORIGIN:\$ORIGIN/.." - INSTALL_RPATH_USE_LINK_PATH ON - SOVERSION ${PROJECT_VERSION_MINOR} - VERSION - ${PROJECT_VERSION_MAJOR}.${PROJECT_VERSION_MINOR}.${PROJECT_VERSION_PATCH}) - -install( - TARGETS kernel-tracing-test-tool - DESTINATION lib/rocprofiler-sdk - COMPONENT tests) - if(ROCPROFILER_MEMCHECK_PRELOAD_ENV) set(PRELOAD_ENV - "${ROCPROFILER_MEMCHECK_PRELOAD_ENV}:$") + "${ROCPROFILER_MEMCHECK_PRELOAD_ENV}:$") else() - set(PRELOAD_ENV "LD_PRELOAD=$") + set(PRELOAD_ENV "LD_PRELOAD=$") endif() add_test(NAME test-kernel-tracing-execute COMMAND $) @@ -46,7 +27,7 @@ set_tests_properties( LABELS "integration-tests" ENVIRONMENT - "${PRELOAD_ENV};HSA_TOOLS_LIB=$" + "${PRELOAD_ENV};HSA_TOOLS_LIB=$;ROCPROFILER_TOOL_OUTPUT_FILE=kernel-tracing-test.json" FAIL_REGULAR_EXPRESSION "threw an exception") @@ -57,7 +38,7 @@ endforeach() add_test(NAME test-kernel-tracing-validate COMMAND ${Python3_EXECUTABLE} ${CMAKE_CURRENT_BINARY_DIR}/validate.py --input - ${CMAKE_CURRENT_BINARY_DIR}/kernel-tracing-test-tool.json) + ${CMAKE_CURRENT_BINARY_DIR}/kernel-tracing-test.json) set_tests_properties( test-kernel-tracing-validate diff --git a/projects/rocprofiler-sdk/tests/kernel-tracing/conftest.py b/projects/rocprofiler-sdk/tests/kernel-tracing/conftest.py index 7a49a1c53b..406128800b 100644 --- a/projects/rocprofiler-sdk/tests/kernel-tracing/conftest.py +++ b/projects/rocprofiler-sdk/tests/kernel-tracing/conftest.py @@ -8,7 +8,7 @@ def pytest_addoption(parser): parser.addoption( "--input", action="store", - default="kernel-tracing-test-tool.json", + default="kernel-tracing-test.json", help="Input JSON", ) diff --git a/projects/rocprofiler-sdk/tests/kernel-tracing/validate.py b/projects/rocprofiler-sdk/tests/kernel-tracing/validate.py index 299359a43d..de27cb1a12 100644 --- a/projects/rocprofiler-sdk/tests/kernel-tracing/validate.py +++ b/projects/rocprofiler-sdk/tests/kernel-tracing/validate.py @@ -15,21 +15,21 @@ def test_data_structure(input_data): """verify minimum amount of expected data is present""" data = input_data - node_exists("kernel-tracing-test-tool", data) - node_exists("agents", data["kernel-tracing-test-tool"]) - node_exists("call_stack", data["kernel-tracing-test-tool"]) - node_exists("callback_records", data["kernel-tracing-test-tool"]) - node_exists("buffer_records", data["kernel-tracing-test-tool"]) + node_exists("rocprofiler-sdk-json-tool", data) + node_exists("agents", data["rocprofiler-sdk-json-tool"]) + node_exists("call_stack", data["rocprofiler-sdk-json-tool"]) + node_exists("callback_records", data["rocprofiler-sdk-json-tool"]) + node_exists("buffer_records", data["rocprofiler-sdk-json-tool"]) - node_exists("names", data["kernel-tracing-test-tool"]["callback_records"]) - node_exists("code_objects", data["kernel-tracing-test-tool"]["callback_records"]) - node_exists("kernel_symbols", data["kernel-tracing-test-tool"]["callback_records"]) - node_exists("hsa_api_traces", data["kernel-tracing-test-tool"]["callback_records"]) + node_exists("names", data["rocprofiler-sdk-json-tool"]["callback_records"]) + node_exists("code_objects", data["rocprofiler-sdk-json-tool"]["callback_records"]) + node_exists("kernel_symbols", data["rocprofiler-sdk-json-tool"]["callback_records"]) + node_exists("hsa_api_traces", data["rocprofiler-sdk-json-tool"]["callback_records"]) - node_exists("names", data["kernel-tracing-test-tool"]["buffer_records"]) - node_exists("kernel_dispatches", data["kernel-tracing-test-tool"]["buffer_records"]) - node_exists("memory_copies", data["kernel-tracing-test-tool"]["buffer_records"], 0) - node_exists("hsa_api_traces", data["kernel-tracing-test-tool"]["buffer_records"]) + node_exists("names", data["rocprofiler-sdk-json-tool"]["buffer_records"]) + node_exists("kernel_dispatches", data["rocprofiler-sdk-json-tool"]["buffer_records"]) + node_exists("memory_copies", data["rocprofiler-sdk-json-tool"]["buffer_records"], 0) + node_exists("hsa_api_traces", data["rocprofiler-sdk-json-tool"]["buffer_records"]) def test_timestamps(input_data): @@ -37,7 +37,7 @@ def test_timestamps(input_data): cb_start = {} cb_end = {} - for itr in data["kernel-tracing-test-tool"]["callback_records"]["hsa_api_traces"]: + for itr in data["rocprofiler-sdk-json-tool"]["callback_records"]["hsa_api_traces"]: cid = itr["record"]["correlation_id"]["internal"] phase = itr["record"]["phase"] if phase == 1: @@ -48,10 +48,13 @@ def test_timestamps(input_data): else: assert phase == 1 or phase == 2 - for itr in data["kernel-tracing-test-tool"]["buffer_records"]["hsa_api_traces"]: + for itr in data["rocprofiler-sdk-json-tool"]["buffer_records"]["hsa_api_traces"]: assert itr["start_timestamp"] <= itr["end_timestamp"] - for itr in data["kernel-tracing-test-tool"]["buffer_records"]["kernel_dispatches"]: + for itr in data["rocprofiler-sdk-json-tool"]["buffer_records"]["memory_copies"]: + assert itr["start_timestamp"] <= itr["end_timestamp"] + + for itr in data["rocprofiler-sdk-json-tool"]["buffer_records"]["kernel_dispatches"]: assert itr["start_timestamp"] < itr["end_timestamp"] assert itr["correlation_id"]["internal"] > 0 assert itr["correlation_id"]["external"] > 0 @@ -66,15 +69,21 @@ def test_internal_correlation_ids(input_data): data = input_data api_corr_ids = [] - for itr in data["kernel-tracing-test-tool"]["callback_records"]["hsa_api_traces"]: + for itr in data["rocprofiler-sdk-json-tool"]["callback_records"]["hsa_api_traces"]: api_corr_ids.append(itr["record"]["correlation_id"]["internal"]) - for itr in data["kernel-tracing-test-tool"]["buffer_records"]["hsa_api_traces"]: + for itr in data["rocprofiler-sdk-json-tool"]["buffer_records"]["hsa_api_traces"]: api_corr_ids.append(itr["correlation_id"]["internal"]) api_corr_ids_sorted = sorted(api_corr_ids) api_corr_ids_unique = list(set(api_corr_ids)) + for itr in data["rocprofiler-sdk-json-tool"]["buffer_records"]["kernel_dispatches"]: + assert itr["correlation_id"]["internal"] in api_corr_ids_unique + + for itr in data["rocprofiler-sdk-json-tool"]["buffer_records"]["memory_copies"]: + assert itr["correlation_id"]["internal"] in api_corr_ids_unique + len_corr_id_unq = len(api_corr_ids_unique) assert len(api_corr_ids) != len_corr_id_unq assert max(api_corr_ids_sorted) == len_corr_id_unq @@ -83,18 +92,33 @@ def test_internal_correlation_ids(input_data): def test_external_correlation_ids(input_data): data = input_data - for itr in data["kernel-tracing-test-tool"]["callback_records"]["hsa_api_traces"]: + extern_corr_ids = [] + for itr in data["rocprofiler-sdk-json-tool"]["callback_records"]["hsa_api_traces"]: + assert itr["record"]["correlation_id"]["external"] > 0 assert itr["record"]["thread_id"] == itr["record"]["correlation_id"]["external"] + extern_corr_ids.append(itr["record"]["correlation_id"]["external"]) - for itr in data["kernel-tracing-test-tool"]["buffer_records"]["hsa_api_traces"]: + extern_corr_ids = list(set(sorted(extern_corr_ids))) + for itr in data["rocprofiler-sdk-json-tool"]["buffer_records"]["hsa_api_traces"]: + assert itr["correlation_id"]["external"] > 0 assert itr["thread_id"] == itr["correlation_id"]["external"] + assert itr["thread_id"] in extern_corr_ids + assert itr["correlation_id"]["external"] in extern_corr_ids + + for itr in data["rocprofiler-sdk-json-tool"]["buffer_records"]["kernel_dispatches"]: + assert itr["correlation_id"]["external"] > 0 + assert itr["correlation_id"]["external"] in extern_corr_ids + + for itr in data["rocprofiler-sdk-json-tool"]["buffer_records"]["memory_copies"]: + assert itr["correlation_id"]["external"] > 0 + assert itr["correlation_id"]["external"] in extern_corr_ids def test_kernel_ids(input_data): data = input_data symbol_info = {} - for itr in data["kernel-tracing-test-tool"]["callback_records"]["kernel_symbols"]: + for itr in data["rocprofiler-sdk-json-tool"]["callback_records"]["kernel_symbols"]: phase = itr["record"]["phase"] payload = itr["payload"] kern_id = payload["kernel_id"] @@ -108,10 +132,33 @@ def test_kernel_ids(input_data): assert payload["kernel_id"] in symbol_info.keys() assert payload["kernel_name"] == symbol_info[kern_id]["kernel_name"] - for itr in data["kernel-tracing-test-tool"]["buffer_records"]["kernel_dispatches"]: + for itr in data["rocprofiler-sdk-json-tool"]["buffer_records"]["kernel_dispatches"]: assert itr["kernel_id"] in symbol_info.keys() +def test_async_copy_direction(input_data): + data = input_data + + # Direction values: + # -1 == ??? (unknown) + # 0 == H2H (host to host) + # 1 == H2D (host to device) + # 2 == D2H (device to host) + # 3 == D2D (device to device) + async_dir_cnt = dict([(idx, 0) for idx in range(-1, 4)]) + for itr in data["rocprofiler-sdk-json-tool"]["buffer_records"]["memory_copies"]: + op_id = itr["operation"] + async_dir_cnt[op_id] += 1 + + # in the reproducible-runtime test which generates the input file, + # we don't expect any async memory copy operations + assert async_dir_cnt[-1] == 0 + assert async_dir_cnt[0] == 0 + assert async_dir_cnt[1] == 0 + assert async_dir_cnt[2] == 0 + assert async_dir_cnt[3] == 0 + + if __name__ == "__main__": exit_code = pytest.main(["-x", __file__] + sys.argv[1:]) sys.exit(exit_code) diff --git a/projects/rocprofiler-sdk/tests/tools/CMakeLists.txt b/projects/rocprofiler-sdk/tests/tools/CMakeLists.txt new file mode 100644 index 0000000000..b1890890d4 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/tools/CMakeLists.txt @@ -0,0 +1,37 @@ +# +# +# +cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR) + +project( + rocprofiler-tests-json-tool + LANGUAGES CXX + VERSION 0.0.0) + +find_package(rocprofiler-sdk REQUIRED) + +add_library(rocprofiler-sdk-json-tool SHARED) +target_sources(rocprofiler-sdk-json-tool PRIVATE json-tool.cpp) +target_link_libraries( + rocprofiler-sdk-json-tool + PRIVATE rocprofiler::rocprofiler rocprofiler::cereal rocprofiler::tests-build-flags + rocprofiler::tests-common-library) +set_target_properties( + rocprofiler-sdk-json-tool + PROPERTIES INSTALL_RPATH "\$ORIGIN:\$ORIGIN/.." + INSTALL_RPATH_USE_LINK_PATH ON + SOVERSION ${PROJECT_VERSION_MINOR} + VERSION + ${PROJECT_VERSION_MAJOR}.${PROJECT_VERSION_MINOR}.${PROJECT_VERSION_PATCH}) + +install( + TARGETS rocprofiler-sdk-json-tool + DESTINATION lib/rocprofiler-sdk + COMPONENT tests) + +if(ROCPROFILER_MEMCHECK_PRELOAD_ENV) + set(PRELOAD_ENV + "${ROCPROFILER_MEMCHECK_PRELOAD_ENV}:$") +else() + set(PRELOAD_ENV "LD_PRELOAD=$") +endif() diff --git a/projects/rocprofiler-sdk/tests/kernel-tracing/kernel-tracing.cpp b/projects/rocprofiler-sdk/tests/tools/json-tool.cpp similarity index 98% rename from projects/rocprofiler-sdk/tests/kernel-tracing/kernel-tracing.cpp rename to projects/rocprofiler-sdk/tests/tools/json-tool.cpp index 551cef4e6f..32557b827d 100644 --- a/projects/rocprofiler-sdk/tests/kernel-tracing/kernel-tracing.cpp +++ b/projects/rocprofiler-sdk/tests/tools/json-tool.cpp @@ -26,7 +26,7 @@ #endif /** - * @file tests/kernel-tracing/kernel-tracing.cpp + * @file tests/tools/json-tool.cpp * * @brief Test rocprofiler tool */ @@ -597,7 +597,7 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data) } // environment variable to select which contexts to collect - auto* context_settings_env = getenv("KERNEL_TRACING_CONTEXTS"); + auto* context_settings_env = getenv("ROCPROFILER_TOOL_CONTEXTS"); if(context_settings_env != nullptr && !std::string_view{context_settings_env}.empty()) { auto context_settings = std::string{context_settings_env}; @@ -623,8 +623,8 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data) { auto filename = std::string_view{__FILE__}; auto msg = std::stringstream{}; - msg << "[kernel-tracing][" << filename.substr(filename.find_last_of('/') + 1) << ":" - << __LINE__ << "] invalid specification of KERNEL_TRACING_CONTEXTS ('" + msg << "[rocprofiler-sdk-json-tool][" << filename.substr(filename.find_last_of('/') + 1) + << ":" << __LINE__ << "] invalid specification of ROCPROFILER_TOOL_CONTEXTS ('" << context_settings_env << "'). Valid choices are: " << options.str(); throw std::runtime_error{msg.str()}; } @@ -661,8 +661,8 @@ tool_fini(void* tool_data) _call_stack->emplace_back(source_location{__FUNCTION__, __FILE__, __LINE__, ""}); } - auto ofname = std::string{"kernel-tracing-test-tool.json"}; - if(auto* eofname = getenv("ROCPROFILER_KERNEL_TRACING_OUTPUT_FILE")) ofname = eofname; + auto ofname = std::string{"rocprofiler-tool-results.json"}; + if(auto* eofname = getenv("ROCPROFILER_TOOL_OUTPUT_FILE")) ofname = eofname; std::ostream* ofs = nullptr; auto cleanup = std::function{}; @@ -699,7 +699,7 @@ tool_fini(void* tool_data) auto buffer_name_info = get_buffer_tracing_names(); auto callback_name_info = get_callback_tracing_names(); - json_ar.setNextName("kernel-tracing-test-tool"); + json_ar.setNextName("rocprofiler-sdk-json-tool"); json_ar.startNode(); json_ar(cereal::make_nvp("agents", agents)); @@ -838,7 +838,7 @@ rocprofiler_configure(uint32_t version, if(priority > 0) return nullptr; // set the client name - id->name = "kernel-tracing-test-tool"; + id->name = "rocprofiler-sdk-json-tool"; // store client info client::client_id = id;