diff --git a/cmake/Templates/rocprofiler-sdk-roctx/config.cmake.in b/cmake/Templates/rocprofiler-sdk-roctx/config.cmake.in index 26483d63b4..e46458b51b 100644 --- a/cmake/Templates/rocprofiler-sdk-roctx/config.cmake.in +++ b/cmake/Templates/rocprofiler-sdk-roctx/config.cmake.in @@ -44,7 +44,7 @@ else() # Library dependencies foreach(TARG @PROJECT_BUILD_TARGETS@) - set(TARG @PACKAGE_NAME@::@PACKAGE_NAME@-${TARG}) + set(TARG @PACKAGE_NAME@::${TARG}) if(NOT @PACKAGE_NAME@_FIND_COMPONENTS) list(APPEND @PACKAGE_NAME@_LIBRARIES ${TARG}) target_link_libraries(@PACKAGE_NAME@::@PACKAGE_NAME@ INTERFACE ${TARG}) @@ -53,7 +53,7 @@ else() if(@PACKAGE_NAME@_FIND_COMPONENTS) foreach(COMP ${@PACKAGE_NAME@_FIND_COMPONENTS}) - set(TARG @PACKAGE_NAME@::${COMP}) + set(TARG @PACKAGE_NAME@::@PACKAGE_NAME@-${COMP}) if(TARGET ${TARG}) set(@PACKAGE_NAME@_${COMP}_FOUND 1) list(APPEND @PACKAGE_NAME@_LIBRARIES ${TARG}) diff --git a/source/include/rocprofiler-sdk/buffer_tracing.h b/source/include/rocprofiler-sdk/buffer_tracing.h index 75e4976ee5..9b011f7b7c 100644 --- a/source/include/rocprofiler-sdk/buffer_tracing.h +++ b/source/include/rocprofiler-sdk/buffer_tracing.h @@ -70,11 +70,12 @@ typedef struct { uint64_t size; ///< size of this struct rocprofiler_buffer_tracing_kind_t kind; ///< ::ROCPROFILER_CALLBACK_TRACING_MARKER_API - rocprofiler_correlation_id_t correlation_id; ///< correlation ids for record - 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 + rocprofiler_correlation_id_t correlation_id; ///< correlation ids for record + rocprofiler_tracing_operation_t operation; ///< ::rocprofiler_marker_api_id_t + rocprofiler_timestamp_t start_timestamp; ///< start time in nanoseconds + rocprofiler_timestamp_t end_timestamp; ///< end time in nanoseconds + rocprofiler_thread_id_t thread_id; ///< id for thread generating this record + uint64_t marker_id; ///< rocprofiler_marker_id_t // const char* message; // (Need Review?) } rocprofiler_buffer_tracing_marker_api_record_t; diff --git a/source/include/rocprofiler-sdk/marker.h b/source/include/rocprofiler-sdk/marker.h index 2bd441a0c9..588e827dc4 100644 --- a/source/include/rocprofiler-sdk/marker.h +++ b/source/include/rocprofiler-sdk/marker.h @@ -24,3 +24,4 @@ #include #include +#include diff --git a/source/include/rocprofiler-sdk/marker/CMakeLists.txt b/source/include/rocprofiler-sdk/marker/CMakeLists.txt index 1833198590..cbf6dae0c2 100644 --- a/source/include/rocprofiler-sdk/marker/CMakeLists.txt +++ b/source/include/rocprofiler-sdk/marker/CMakeLists.txt @@ -3,7 +3,7 @@ # Installation of public HSA headers # # -set(ROCPROFILER_MARKER_HEADER_FILES api_args.h api_id.h) +set(ROCPROFILER_MARKER_HEADER_FILES api_args.h api_id.h table_api_id.h) install( FILES ${ROCPROFILER_MARKER_HEADER_FILES} diff --git a/source/include/rocprofiler-sdk/marker/api_args.h b/source/include/rocprofiler-sdk/marker/api_args.h index 9a556dbb8f..a07c688993 100644 --- a/source/include/rocprofiler-sdk/marker/api_args.h +++ b/source/include/rocprofiler-sdk/marker/api_args.h @@ -22,13 +22,15 @@ #pragma once -#include +#include +#include -typedef uint64_t roctx_range_id_t; +#include typedef union rocprofiler_marker_api_retval_u { - uint32_t uint32_t_retval; + int32_t int32_t_retval; + int64_t int64_t_retval; uint64_t uint64_t_retval; } rocprofiler_marker_api_retval_t; @@ -53,4 +55,35 @@ typedef union rocprofiler_marker_api_args_u { roctx_range_id_t id; } roctxRangeStop; + struct + { + roctx_thread_id_t tid; + } roctxProfilerPause; + struct + { + roctx_thread_id_t tid; + } roctxProfilerResume; + struct + { + const char* name; + } roctxNameOsThread; + struct + { + const char* name; + const struct hsa_agent_s* agent; + } roctxNameHsaAgent; + struct + { + const char* name; + int device_id; + } roctxNameHipDevice; + struct + { + const char* name; + const struct ihipStream_t* stream; + } roctxNameHipStream; + struct + { + roctx_thread_id_t* tid; + } roctxGetThreadId; } rocprofiler_marker_api_args_t; diff --git a/source/include/rocprofiler-sdk/marker/api_id.h b/source/include/rocprofiler-sdk/marker/api_id.h index dddd3b8cb4..b6f7c51179 100644 --- a/source/include/rocprofiler-sdk/marker/api_id.h +++ b/source/include/rocprofiler-sdk/marker/api_id.h @@ -36,8 +36,9 @@ typedef enum // NOLINT(performance-enum-size) ROCPROFILER_MARKER_API_ID_roctxProfilerPause, ROCPROFILER_MARKER_API_ID_roctxProfilerResume, ROCPROFILER_MARKER_API_ID_roctxNameOsThread, - ROCPROFILER_MARKER_API_ID_roctxNameHsaDevice, + ROCPROFILER_MARKER_API_ID_roctxNameHsaAgent, ROCPROFILER_MARKER_API_ID_roctxNameHipDevice, ROCPROFILER_MARKER_API_ID_roctxNameHipStream, + ROCPROFILER_MARKER_API_ID_roctxGetThreadId, ROCPROFILER_MARKER_API_ID_LAST, } rocprofiler_marker_api_id_t; diff --git a/source/include/rocprofiler-sdk/marker/table_api_id.h b/source/include/rocprofiler-sdk/marker/table_api_id.h new file mode 100644 index 0000000000..ceed829699 --- /dev/null +++ b/source/include/rocprofiler-sdk/marker/table_api_id.h @@ -0,0 +1,30 @@ +// 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 + +// NOLINTNEXTLINE(performance-enum-size) +typedef enum +{ + ROCPROFILER_MARKER_API_TABLE_ID_NONE = -1, + ROCPROFILER_MARKER_API_TABLE_ID_RoctxApi = 0, +} rocprofiler_marker_table_api_id_t; diff --git a/source/lib/common/logging.cpp b/source/lib/common/logging.cpp index 56d37349e7..3d4938ef8c 100644 --- a/source/lib/common/logging.cpp +++ b/source/lib/common/logging.cpp @@ -79,7 +79,7 @@ init_logging(std::string_view env_var) FLAGS_minloglevel = loglvl_v; FLAGS_stderrthreshold = loglvl_v; - LOG(INFO) << "logging initialized"; + LOG(INFO) << "logging initialized via " << env_var; }); } } // namespace common diff --git a/source/lib/common/mpl.hpp b/source/lib/common/mpl.hpp index 1f05912ad9..f37180028d 100644 --- a/source/lib/common/mpl.hpp +++ b/source/lib/common/mpl.hpp @@ -90,6 +90,28 @@ struct is_one_of> { static constexpr bool value = is_one_of>::value; }; + +template +struct is_pair_impl +{ + static constexpr auto value = false; +}; + +template +struct is_pair_impl> +{ + static constexpr auto value = true; +}; + +template +struct is_pair : is_pair_impl>>> +{}; + +template +constexpr bool is_type_complete_v = false; // NOLINT(misc-definitions-in-headers) + +template // NOLINTNEXTLINE(misc-definitions-in-headers) +constexpr bool is_type_complete_v> = true; } // namespace mpl } // namespace common } // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk-roctx/CMakeLists.txt b/source/lib/rocprofiler-sdk-roctx/CMakeLists.txt index 15354db9c4..ef51dd883a 100644 --- a/source/lib/rocprofiler-sdk-roctx/CMakeLists.txt +++ b/source/lib/rocprofiler-sdk-roctx/CMakeLists.txt @@ -7,6 +7,8 @@ find_package(rocprofiler-register REQUIRED) add_library(rocprofiler-sdk-roctx-shared-library SHARED) add_library(rocprofiler::rocprofiler-sdk-roctx-shared-library ALIAS rocprofiler-sdk-roctx-shared-library) +add_library(rocprofiler-sdk-roctx::rocprofiler-sdk-roctx-shared-library ALIAS + rocprofiler-sdk-roctx-shared-library) target_sources(rocprofiler-sdk-roctx-shared-library PRIVATE roctx.cpp) target_include_directories( diff --git a/source/lib/rocprofiler-sdk/CMakeLists.txt b/source/lib/rocprofiler-sdk/CMakeLists.txt index a168d5a8e2..17b9d25d7a 100644 --- a/source/lib/rocprofiler-sdk/CMakeLists.txt +++ b/source/lib/rocprofiler-sdk/CMakeLists.txt @@ -40,6 +40,7 @@ add_subdirectory(context) add_subdirectory(counters) add_subdirectory(aql) add_subdirectory(pc_sampling) +add_subdirectory(marker) target_link_libraries( rocprofiler-object-library diff --git a/source/lib/rocprofiler-sdk/agent.cpp b/source/lib/rocprofiler-sdk/agent.cpp index 7a39f1b6b6..f873cc9545 100644 --- a/source/lib/rocprofiler-sdk/agent.cpp +++ b/source/lib/rocprofiler-sdk/agent.cpp @@ -30,7 +30,7 @@ #include "lib/rocprofiler-sdk/agent.hpp" #include "lib/rocprofiler-sdk/hsa/agent_cache.hpp" -#include +#include #include #include #include @@ -404,7 +404,8 @@ read_topology() agent_info.node_id = nodecount++; if(!name_prop.empty()) - agent_info.model_name = get_string_entry(name_prop.front())->c_str(); + agent_info.model_name = + get_string_entry(fmt::format("{}", fmt::join(name_prop, " ")))->c_str(); else agent_info.model_name = ""; diff --git a/source/lib/rocprofiler-sdk/buffer_tracing.cpp b/source/lib/rocprofiler-sdk/buffer_tracing.cpp index cc59cfbd22..063ca8db49 100644 --- a/source/lib/rocprofiler-sdk/buffer_tracing.cpp +++ b/source/lib/rocprofiler-sdk/buffer_tracing.cpp @@ -26,6 +26,7 @@ #include "lib/rocprofiler-sdk/context/context.hpp" #include "lib/rocprofiler-sdk/context/domain.hpp" #include "lib/rocprofiler-sdk/hsa/hsa.hpp" +#include "lib/rocprofiler-sdk/marker/marker.hpp" #include "lib/rocprofiler-sdk/registration.hpp" #include @@ -143,25 +144,26 @@ rocprofiler_query_buffer_tracing_kind_operation_name(rocprofiler_buffer_tracing_ if(kind < ROCPROFILER_BUFFER_TRACING_NONE || kind >= ROCPROFILER_BUFFER_TRACING_LAST) return ROCPROFILER_STATUS_ERROR_KIND_NOT_FOUND; + const char* val = nullptr; if(kind == ROCPROFILER_BUFFER_TRACING_HSA_API) + val = rocprofiler::hsa::name_by_id(operation); + else if(kind == ROCPROFILER_BUFFER_TRACING_MARKER_API) + val = rocprofiler::marker::name_by_id(operation); + else + return ROCPROFILER_STATUS_ERROR_NOT_IMPLEMENTED; + + if(!val) { - const auto* val = rocprofiler::hsa::name_by_id(operation); + if(name) *name = nullptr; + if(name_len) *name_len = 0; - if(!val) - { - if(name) *name = nullptr; - if(name_len) *name_len = 0; - - return ROCPROFILER_STATUS_ERROR_OPERATION_NOT_FOUND; - } - - if(name) *name = val; - if(name_len) *name_len = strnlen(val, 4096); - - return ROCPROFILER_STATUS_SUCCESS; + return ROCPROFILER_STATUS_ERROR_OPERATION_NOT_FOUND; } - return ROCPROFILER_STATUS_ERROR_NOT_IMPLEMENTED; + if(name) *name = val; + if(name_len) *name_len = strnlen(val, 4096); + + return ROCPROFILER_STATUS_SUCCESS; } rocprofiler_status_t @@ -182,18 +184,20 @@ rocprofiler_iterate_buffer_tracing_kind_operations( rocprofiler_buffer_tracing_kind_operation_cb_t callback, void* data) { + auto ops = std::vector{}; if(kind == ROCPROFILER_BUFFER_TRACING_HSA_API) - { - auto ops = rocprofiler::hsa::get_ids(); - for(const auto& itr : ops) - { - auto _success = callback(kind, itr, data); - if(_success != 0) break; - } - return ROCPROFILER_STATUS_SUCCESS; - } + ops = rocprofiler::hsa::get_ids(); + else if(kind == ROCPROFILER_BUFFER_TRACING_MARKER_API) + ops = rocprofiler::marker::get_ids(); + else + return ROCPROFILER_STATUS_ERROR_NOT_IMPLEMENTED; - return ROCPROFILER_STATUS_ERROR_NOT_IMPLEMENTED; + for(const auto& itr : ops) + { + auto _success = callback(kind, itr, data); + if(_success != 0) break; + } + return ROCPROFILER_STATUS_SUCCESS; } } diff --git a/source/lib/rocprofiler-sdk/callback_tracing.cpp b/source/lib/rocprofiler-sdk/callback_tracing.cpp index 8dea70a86f..d505dd8537 100644 --- a/source/lib/rocprofiler-sdk/callback_tracing.cpp +++ b/source/lib/rocprofiler-sdk/callback_tracing.cpp @@ -26,6 +26,7 @@ #include "lib/rocprofiler-sdk/context/context.hpp" #include "lib/rocprofiler-sdk/context/domain.hpp" #include "lib/rocprofiler-sdk/hsa/hsa.hpp" +#include "lib/rocprofiler-sdk/marker/marker.hpp" #include "lib/rocprofiler-sdk/registration.hpp" #include @@ -134,25 +135,26 @@ rocprofiler_query_callback_tracing_kind_operation_name(rocprofiler_callback_trac if(kind < ROCPROFILER_CALLBACK_TRACING_NONE || kind >= ROCPROFILER_CALLBACK_TRACING_LAST) return ROCPROFILER_STATUS_ERROR_KIND_NOT_FOUND; + const char* val = nullptr; if(kind == ROCPROFILER_CALLBACK_TRACING_HSA_API) + val = rocprofiler::hsa::name_by_id(operation); + else if(kind == ROCPROFILER_CALLBACK_TRACING_MARKER_API) + val = rocprofiler::marker::name_by_id(operation); + else + return ROCPROFILER_STATUS_ERROR_NOT_IMPLEMENTED; + + if(!val) { - const auto* val = rocprofiler::hsa::name_by_id(operation); + if(name) *name = nullptr; + if(name_len) *name_len = 0; - if(!val) - { - if(name) *name = nullptr; - if(name_len) *name_len = 0; - - return ROCPROFILER_STATUS_ERROR_OPERATION_NOT_FOUND; - } - - if(name) *name = val; - if(name_len) *name_len = strnlen(val, 4096); - - return ROCPROFILER_STATUS_SUCCESS; + return ROCPROFILER_STATUS_ERROR_OPERATION_NOT_FOUND; } - return ROCPROFILER_STATUS_ERROR_NOT_IMPLEMENTED; + if(name) *name = val; + if(name_len) *name_len = strnlen(val, 4096); + + return ROCPROFILER_STATUS_SUCCESS; } rocprofiler_status_t @@ -174,18 +176,20 @@ rocprofiler_iterate_callback_tracing_kind_operations( rocprofiler_callback_tracing_kind_operation_cb_t callback, void* data) { + auto ops = std::vector{}; if(kind == ROCPROFILER_CALLBACK_TRACING_HSA_API) - { - auto ops = rocprofiler::hsa::get_ids(); - for(const auto& itr : ops) - { - auto _success = callback(kind, itr, data); - if(_success != 0) break; - } - return ROCPROFILER_STATUS_SUCCESS; - } + ops = rocprofiler::hsa::get_ids(); + else if(kind == ROCPROFILER_CALLBACK_TRACING_MARKER_API) + ops = rocprofiler::marker::get_ids(); + else + return ROCPROFILER_STATUS_ERROR_NOT_IMPLEMENTED; - return ROCPROFILER_STATUS_ERROR_NOT_IMPLEMENTED; + for(const auto& itr : ops) + { + auto _success = callback(kind, itr, data); + if(_success != 0) break; + } + return ROCPROFILER_STATUS_SUCCESS; } rocprofiler_status_t @@ -203,6 +207,15 @@ rocprofiler_iterate_callback_tracing_kind_operation_args( user_data); return ROCPROFILER_STATUS_SUCCESS; } + else if(record.kind == ROCPROFILER_CALLBACK_TRACING_MARKER_API) + { + rocprofiler::marker::iterate_args( + record.operation, + *static_cast(record.payload), + callback, + user_data); + return ROCPROFILER_STATUS_SUCCESS; + } return ROCPROFILER_STATUS_ERROR_NOT_IMPLEMENTED; } diff --git a/source/lib/rocprofiler-sdk/intercept_table.cpp b/source/lib/rocprofiler-sdk/intercept_table.cpp index 399e20bd6a..c5fa8f2429 100644 --- a/source/lib/rocprofiler-sdk/intercept_table.cpp +++ b/source/lib/rocprofiler-sdk/intercept_table.cpp @@ -156,6 +156,11 @@ template void notify_runtime_api_registration(rocprofiler_runtime_library_t, uint64_t, uint64_t, std::tuple); + +template void notify_runtime_api_registration(rocprofiler_runtime_library_t, + uint64_t, + uint64_t, + std::tuple); } // namespace intercept_table } // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk/marker/CMakeLists.txt b/source/lib/rocprofiler-sdk/marker/CMakeLists.txt new file mode 100644 index 0000000000..68a87f38f3 --- /dev/null +++ b/source/lib/rocprofiler-sdk/marker/CMakeLists.txt @@ -0,0 +1,7 @@ +# +# +set(ROCPROFILER_LIB_MARKER_SOURCES marker.cpp) +set(ROCPROFILER_LIB_MARKER_HEADERS defines.hpp marker.hpp utils.hpp) + +target_sources(rocprofiler-object-library PRIVATE ${ROCPROFILER_LIB_MARKER_SOURCES} + ${ROCPROFILER_LIB_MARKER_HEADERS}) diff --git a/source/lib/rocprofiler-sdk/marker/defines.hpp b/source/lib/rocprofiler-sdk/marker/defines.hpp new file mode 100644 index 0000000000..7f0d8c980b --- /dev/null +++ b/source/lib/rocprofiler-sdk/marker/defines.hpp @@ -0,0 +1,286 @@ +// 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 + +#define IMPL_DETAIL_EXPAND(X) X +#define IMPL_DETAIL_FOR_EACH_NARG(...) \ + IMPL_DETAIL_FOR_EACH_NARG_(__VA_ARGS__, IMPL_DETAIL_FOR_EACH_RSEQ_N()) +#define IMPL_DETAIL_FOR_EACH_NARG_(...) IMPL_DETAIL_EXPAND(IMPL_DETAIL_FOR_EACH_ARG_N(__VA_ARGS__)) +#define IMPL_DETAIL_FOR_EACH_ARG_N( \ + _1, _2, _3, _4, _5, _6, _7, _8, _9, _10, _11, _12, _13, _14, _15, N, ...) \ + N +#define IMPL_DETAIL_FOR_EACH_RSEQ_N() 15, 14, 13, 12, 11, 10, 9, 8, 7, 6, 5, 4, 3, 2, 1, 0 +#define IMPL_DETAIL_CONCATENATE(X, Y) X##Y +#define IMPL_DETAIL_FOR_EACH_(N, MACRO, PREFIX, ...) \ + IMPL_DETAIL_EXPAND(IMPL_DETAIL_CONCATENATE(MACRO, N)(PREFIX, __VA_ARGS__)) +#define IMPL_DETAIL_FOR_EACH(MACRO, PREFIX, ...) \ + IMPL_DETAIL_FOR_EACH_(IMPL_DETAIL_FOR_EACH_NARG(__VA_ARGS__), MACRO, PREFIX, __VA_ARGS__) + +#define ADDR_MEMBER_0(...) +#define ADDR_MEMBER_1(PREFIX, FIELD) static_cast(&PREFIX.FIELD) +#define ADDR_MEMBER_2(PREFIX, A, B) ADDR_MEMBER_1(PREFIX, A), ADDR_MEMBER_1(PREFIX, B) +#define ADDR_MEMBER_3(PREFIX, A, B, C) ADDR_MEMBER_2(PREFIX, A, B), ADDR_MEMBER_1(PREFIX, C) +#define ADDR_MEMBER_4(PREFIX, A, B, C, D) ADDR_MEMBER_3(PREFIX, A, B, C), ADDR_MEMBER_1(PREFIX, D) +#define ADDR_MEMBER_5(PREFIX, A, B, C, D, E) \ + ADDR_MEMBER_4(PREFIX, A, B, C, D), ADDR_MEMBER_1(PREFIX, E) +#define ADDR_MEMBER_6(PREFIX, A, B, C, D, E, F) \ + ADDR_MEMBER_5(PREFIX, A, B, C, D, E), ADDR_MEMBER_1(PREFIX, F) +#define ADDR_MEMBER_7(PREFIX, A, B, C, D, E, F, G) \ + ADDR_MEMBER_6(PREFIX, A, B, C, D, E, F), ADDR_MEMBER_1(PREFIX, G) +#define ADDR_MEMBER_8(PREFIX, A, B, C, D, E, F, G, H) \ + ADDR_MEMBER_7(PREFIX, A, B, C, D, E, F, G), ADDR_MEMBER_1(PREFIX, H) +#define ADDR_MEMBER_9(PREFIX, A, B, C, D, E, F, G, H, I) \ + ADDR_MEMBER_8(PREFIX, A, B, C, D, E, F, G, H), ADDR_MEMBER_1(PREFIX, I) +#define ADDR_MEMBER_10(PREFIX, A, B, C, D, E, F, G, H, I, J) \ + ADDR_MEMBER_9(PREFIX, A, B, C, D, E, F, G, H, I), ADDR_MEMBER_1(PREFIX, J) +#define ADDR_MEMBER_11(PREFIX, A, B, C, D, E, F, G, H, I, J, K) \ + ADDR_MEMBER_10(PREFIX, A, B, C, D, E, F, G, H, I, J), ADDR_MEMBER_1(PREFIX, K) +#define ADDR_MEMBER_12(PREFIX, A, B, C, D, E, F, G, H, I, J, K, L) \ + ADDR_MEMBER_11(PREFIX, A, B, C, D, E, F, G, H, I, J, K), ADDR_MEMBER_1(PREFIX, L) +#define ADDR_MEMBER_13(PREFIX, A, B, C, D, E, F, G, H, I, J, K, L, M) \ + ADDR_MEMBER_12(PREFIX, A, B, C, D, E, F, G, H, I, J, K, L), ADDR_MEMBER_1(PREFIX, M) +#define ADDR_MEMBER_14(PREFIX, A, B, C, D, E, F, G, H, I, J, K, L, M, N) \ + ADDR_MEMBER_13(PREFIX, A, B, C, D, E, F, G, H, I, J, K, L, M), ADDR_MEMBER_1(PREFIX, N) +#define ADDR_MEMBER_15(PREFIX, A, B, C, D, E, F, G, H, I, J, K, L, M, N, O) \ + ADDR_MEMBER_14(PREFIX, A, B, C, D, E, F, G, H, I, J, K, L, M, N), ADDR_MEMBER_1(PREFIX, O) + +#define NAMED_MEMBER_0(...) +#define NAMED_MEMBER_1(PREFIX, FIELD) std::make_pair(#FIELD, PREFIX.FIELD) +#define NAMED_MEMBER_2(PREFIX, A, B) NAMED_MEMBER_1(PREFIX, A), NAMED_MEMBER_1(PREFIX, B) +#define NAMED_MEMBER_3(PREFIX, A, B, C) NAMED_MEMBER_2(PREFIX, A, B), NAMED_MEMBER_1(PREFIX, C) +#define NAMED_MEMBER_4(PREFIX, A, B, C, D) \ + NAMED_MEMBER_3(PREFIX, A, B, C), NAMED_MEMBER_1(PREFIX, D) +#define NAMED_MEMBER_5(PREFIX, A, B, C, D, E) \ + NAMED_MEMBER_4(PREFIX, A, B, C, D), NAMED_MEMBER_1(PREFIX, E) +#define NAMED_MEMBER_6(PREFIX, A, B, C, D, E, F) \ + NAMED_MEMBER_5(PREFIX, A, B, C, D, E), NAMED_MEMBER_1(PREFIX, F) +#define NAMED_MEMBER_7(PREFIX, A, B, C, D, E, F, G) \ + NAMED_MEMBER_6(PREFIX, A, B, C, D, E, F), NAMED_MEMBER_1(PREFIX, G) +#define NAMED_MEMBER_8(PREFIX, A, B, C, D, E, F, G, H) \ + NAMED_MEMBER_7(PREFIX, A, B, C, D, E, F, G), NAMED_MEMBER_1(PREFIX, H) +#define NAMED_MEMBER_9(PREFIX, A, B, C, D, E, F, G, H, I) \ + NAMED_MEMBER_8(PREFIX, A, B, C, D, E, F, G, H), NAMED_MEMBER_1(PREFIX, I) +#define NAMED_MEMBER_10(PREFIX, A, B, C, D, E, F, G, H, I, J) \ + NAMED_MEMBER_9(PREFIX, A, B, C, D, E, F, G, H, I), NAMED_MEMBER_1(PREFIX, J) +#define NAMED_MEMBER_11(PREFIX, A, B, C, D, E, F, G, H, I, J, K) \ + NAMED_MEMBER_10(PREFIX, A, B, C, D, E, F, G, H, I, J), NAMED_MEMBER_1(PREFIX, K) +#define NAMED_MEMBER_12(PREFIX, A, B, C, D, E, F, G, H, I, J, K, L) \ + NAMED_MEMBER_11(PREFIX, A, B, C, D, E, F, G, H, I, J, K), NAMED_MEMBER_1(PREFIX, L) +#define NAMED_MEMBER_13(PREFIX, A, B, C, D, E, F, G, H, I, J, K, L, M) \ + NAMED_MEMBER_12(PREFIX, A, B, C, D, E, F, G, H, I, J, K, L), NAMED_MEMBER_1(PREFIX, M) +#define NAMED_MEMBER_14(PREFIX, A, B, C, D, E, F, G, H, I, J, K, L, M, N) \ + NAMED_MEMBER_13(PREFIX, A, B, C, D, E, F, G, H, I, J, K, L, M), NAMED_MEMBER_1(PREFIX, N) +#define NAMED_MEMBER_15(PREFIX, A, B, C, D, E, F, G, H, I, J, K, L, M, N, O) \ + NAMED_MEMBER_14(PREFIX, A, B, C, D, E, F, G, H, I, J, K, L, M, N), NAMED_MEMBER_1(PREFIX, O) + +#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 MARKER_API_INFO_DEFINITION_0(MARKER_TABLE, MARKER_API_ID, MARKER_FUNC, MARKER_FUNC_PTR) \ + namespace rocprofiler \ + { \ + namespace marker \ + { \ + template <> \ + struct roctx_api_info : roctx_domain_info \ + { \ + static constexpr auto table_idx = MARKER_TABLE; \ + static constexpr auto operation_idx = MARKER_API_ID; \ + static constexpr auto name = #MARKER_FUNC; \ + \ + using domain_type = roctx_domain_info; \ + using this_type = roctx_api_info; \ + using base_type = roctx_api_impl; \ + \ + using domain_type::callback_domain_idx; \ + using domain_type::buffered_domain_idx; \ + using domain_type::args_type; \ + using domain_type::retval_type; \ + using domain_type::callback_data_type; \ + \ + static constexpr auto offset() \ + { \ + return offsetof(roctx_table_lookup::type, MARKER_FUNC_PTR); \ + } \ + \ + static_assert(offsetof(roctx_table_lookup::type, MARKER_FUNC_PTR) == \ + (sizeof(size_t) + (operation_idx * sizeof(void*))), \ + "ABI error for " #MARKER_FUNC); \ + \ + static auto& get_table() { return roctx_table_lookup{}(); } \ + \ + template \ + static auto& get_table(TableT& _v) \ + { \ + return roctx_table_lookup{}(_v); \ + } \ + \ + template \ + static auto& get_table_func(TableT& _table) \ + { \ + if constexpr(std::is_pointer::value) \ + { \ + assert(_table != nullptr && "nullptr to MARKER table for " #MARKER_FUNC \ + " function"); \ + return _table->MARKER_FUNC_PTR; \ + } \ + else \ + { \ + return _table.MARKER_FUNC_PTR; \ + } \ + } \ + \ + static auto& get_table_func() { return get_table_func(get_table()); } \ + \ + template \ + static auto& get_api_data_args(DataT& _data) \ + { \ + return _data.MARKER_FUNC; \ + } \ + \ + template \ + static auto get_functor(RetT (*)(Args...)) \ + { \ + if constexpr(std::is_void::value) \ + return [](Args... args) -> RetT { base_type::functor(args...); }; \ + else \ + return [](Args... args) -> RetT { return base_type::functor(args...); }; \ + } \ + \ + static auto get_functor() { return get_functor(get_table_func()); } \ + \ + static std::vector as_arg_addr(callback_data_type) { return std::vector{}; } \ + \ + static std::vector> as_arg_list(callback_data_type) \ + { \ + return {}; \ + } \ + }; \ + } \ + } + +#define MARKER_API_INFO_DEFINITION_V( \ + MARKER_TABLE, MARKER_API_ID, MARKER_FUNC, MARKER_FUNC_PTR, ...) \ + namespace rocprofiler \ + { \ + namespace marker \ + { \ + template <> \ + struct roctx_api_info : roctx_domain_info \ + { \ + static constexpr auto table_idx = MARKER_TABLE; \ + static constexpr auto operation_idx = MARKER_API_ID; \ + static constexpr auto name = #MARKER_FUNC; \ + \ + using domain_type = roctx_domain_info; \ + using this_type = roctx_api_info; \ + using base_type = roctx_api_impl; \ + \ + static constexpr auto callback_domain_idx = domain_type::callback_domain_idx; \ + static constexpr auto buffered_domain_idx = domain_type::buffered_domain_idx; \ + \ + using domain_type::args_type; \ + using domain_type::retval_type; \ + using domain_type::callback_data_type; \ + \ + static constexpr auto offset() \ + { \ + return offsetof(roctx_table_lookup::type, MARKER_FUNC_PTR); \ + } \ + \ + static_assert(offsetof(roctx_table_lookup::type, MARKER_FUNC_PTR) == \ + (sizeof(size_t) + (operation_idx * sizeof(void*))), \ + "ABI error for " #MARKER_FUNC); \ + \ + static auto& get_table() { return roctx_table_lookup{}(); } \ + \ + template \ + static auto& get_table(TableT& _v) \ + { \ + return roctx_table_lookup{}(_v); \ + } \ + \ + template \ + static auto& get_table_func(TableT& _table) \ + { \ + if constexpr(std::is_pointer::value) \ + { \ + assert(_table != nullptr && "nullptr to MARKER table for " #MARKER_FUNC \ + " function"); \ + return _table->MARKER_FUNC_PTR; \ + } \ + else \ + { \ + return _table.MARKER_FUNC_PTR; \ + } \ + } \ + \ + static auto& get_table_func() { return get_table_func(get_table()); } \ + \ + template \ + static auto& get_api_data_args(DataT& _data) \ + { \ + return _data.MARKER_FUNC; \ + } \ + \ + template \ + static auto get_functor(RetT (*)(Args...)) \ + { \ + if constexpr(std::is_same::value) \ + return [](Args... args) -> RetT { base_type::functor(args...); }; \ + else \ + return [](Args... args) -> RetT { return base_type::functor(args...); }; \ + } \ + \ + static auto get_functor() { return get_functor(get_table_func()); } \ + \ + static std::vector as_arg_addr(callback_data_type trace_data) \ + { \ + return std::vector{ \ + GET_ADDR_MEMBER_FIELDS(get_api_data_args(trace_data.args), __VA_ARGS__)}; \ + } \ + \ + static auto as_arg_list(callback_data_type trace_data) \ + { \ + return utils::stringize( \ + GET_NAMED_MEMBER_FIELDS(get_api_data_args(trace_data.args), __VA_ARGS__)); \ + } \ + }; \ + } \ + } + +#define MARKER_API_TABLE_LOOKUP_DEFINITION(TABLE_ID, TYPE) \ + namespace rocprofiler \ + { \ + namespace marker \ + { \ + template <> \ + struct roctx_table_lookup \ + { \ + using type = TYPE; \ + auto& operator()(roctx_api_table_t& _v) const { return _v; } \ + auto& operator()(roctx_api_table_t* _v) const { return *_v; } \ + auto& operator()() const { return (*this)(get_table()); } \ + }; \ + } \ + } diff --git a/source/lib/rocprofiler-sdk/marker/marker.cpp b/source/lib/rocprofiler-sdk/marker/marker.cpp new file mode 100644 index 0000000000..8081051973 --- /dev/null +++ b/source/lib/rocprofiler-sdk/marker/marker.cpp @@ -0,0 +1,612 @@ +// 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/marker/marker.hpp" +#include "lib/common/defines.hpp" +#include "lib/common/utility.hpp" +#include "lib/rocprofiler-sdk/buffer.hpp" +#include "lib/rocprofiler-sdk/context/context.hpp" +#include "lib/rocprofiler-sdk/marker/utils.hpp" + +#include +#include +#include +#include + +#include +#include + +#include +#include +#include +#include +#include + +namespace rocprofiler +{ +namespace marker +{ +namespace +{ +struct null_type +{}; + +template +auto +get_default_retval() +{ + if constexpr(std::is_integral::value) + return Tp{0}; + else + static_assert(std::is_empty::value, "Error! unsupported return type"); +} + +template +void +set_data_retval(DataT& _data, Tp _val) +{ + if constexpr(std::is_same::value) + _data.int32_t_retval = _val; + else if constexpr(std::is_same::value) + _data.int64_t_retval = _val; + else if constexpr(std::is_same::value) + _data.uint64_t_retval = _val; + else + static_assert(std::is_empty::value, "Error! unsupported return type"); +} +} // namespace + +roctx_api_table_t& +get_table() +{ + static auto _v = common::init_public_api_struct(roctx_api_table_t{}); + return _v; +} + +template +template +auto +roctx_api_impl::set_data_args(DataArgsT& _data_args, Args... args) +{ + _data_args = DataArgsT{args...}; +} + +template +template +auto +roctx_api_impl::exec(FuncT&& _func, Args&&... args) +{ + using return_type = std::decay_t>; + + if(_func) + { + if constexpr(std::is_void::value) + { + _func(std::forward(args)...); + return null_type{}; + } + else + { + return _func(std::forward(args)...); + } + } + + using info_type = roctx_api_info; + LOG(ERROR) << "nullptr to next roctx function for " << info_type::name << " (" + << info_type::operation_idx << ")"; + + if constexpr(std::is_void::value) + return null_type{}; + else + return get_default_retval(); +} + +namespace +{ +using correlation_service = context::correlation_tracing_service; +using buffer_marker_api_record_t = rocprofiler_buffer_tracing_marker_api_record_t; + +struct callback_context_data +{ + const context::context* ctx = nullptr; + rocprofiler_callback_tracing_record_t record = {}; + rocprofiler_user_data_t user_data = {.value = 0}; +}; + +struct buffered_context_data +{ + const context::context* ctx = nullptr; + rocprofiler_user_data_t external_correlation = {}; +}; + +constexpr auto empty_user_data = rocprofiler_user_data_t{.value = 0}; + +void +populate_contexts(rocprofiler_callback_tracing_kind_t callback_domain_idx, + rocprofiler_buffer_tracing_kind_t buffered_domain_idx, + int operation_idx, + std::vector& callback_contexts, + std::vector& buffered_contexts) +{ + auto active_contexts = context::context_array_t{}; + auto thr_id = common::get_tid(); + for(const auto* itr : context::get_active_contexts(active_contexts)) + { + if(!itr) continue; + + // if(itr->pc_sampler) has_pc_sampling = true; + + if(itr->callback_tracer) + { + // if the given domain + op is not enabled, skip this context + if(itr->callback_tracer->domains(callback_domain_idx, operation_idx)) + callback_contexts.emplace_back( + callback_context_data{itr, rocprofiler_callback_tracing_record_t{}}); + } + + if(itr->buffered_tracer) + { + // if the given domain + op is not enabled, skip this context + if(itr->buffered_tracer->domains(buffered_domain_idx, operation_idx)) + buffered_contexts.emplace_back(buffered_context_data{ + itr, itr->correlation_tracer.external_correlator.get(thr_id)}); + } + } +} +} // namespace + +template +template +auto +roctx_api_impl::functor(Args&&... args) +{ + using info_type = roctx_api_info; + using callback_api_data_t = typename roctx_domain_info::callback_data_type; + + auto thr_id = common::get_tid(); + auto callback_contexts = std::vector{}; + auto buffered_contexts = std::vector{}; + auto has_pc_sampling = false; + + populate_contexts(info_type::callback_domain_idx, + info_type::buffered_domain_idx, + info_type::operation_idx, + callback_contexts, + buffered_contexts); + + if(callback_contexts.empty() && buffered_contexts.empty()) + { + auto _ret = exec(info_type::get_table_func(), std::forward(args)...); + if constexpr(!std::is_same::value) + return _ret; + else + return 0; + } + + auto ref_count = (has_pc_sampling) ? 4 : 2; + auto buffer_record = common::init_public_api_struct(buffer_marker_api_record_t{}); + auto tracer_data = callback_api_data_t{.size = sizeof(callback_api_data_t)}; + auto* corr_id = correlation_service::construct(ref_count); + auto internal_corr_id = corr_id->internal; + + // construct the buffered info before the callback so the callbacks are as closely wrapped + // around the function call as possible + if(!buffered_contexts.empty()) + { + buffer_record.kind = info_type::buffered_domain_idx; + // external correlation will be updated right before record is placed in buffer + buffer_record.correlation_id = + rocprofiler_correlation_id_t{internal_corr_id, empty_user_data}; + buffer_record.operation = info_type::operation_idx; + buffer_record.thread_id = thr_id; + } + + tracer_data.size = sizeof(callback_api_data_t); + set_data_args(info_type::get_api_data_args(tracer_data.args), std::forward(args)...); + + // invoke the callbacks + if(!callback_contexts.empty()) + { + set_data_args(info_type::get_api_data_args(tracer_data.args), std::forward(args)...); + + for(auto& itr : callback_contexts) + { + auto& ctx = itr.ctx; + auto& record = itr.record; + auto& user_data = itr.user_data; + + auto extern_corr_id_v = ctx->correlation_tracer.external_correlator.get(thr_id); + + auto corr_id_v = rocprofiler_correlation_id_t{internal_corr_id, extern_corr_id_v}; + record = + rocprofiler_callback_tracing_record_t{rocprofiler_context_id_t{ctx->context_idx}, + thr_id, + corr_id_v, + info_type::callback_domain_idx, + info_type::operation_idx, + ROCPROFILER_CALLBACK_PHASE_ENTER, + static_cast(&tracer_data)}; + + auto& callback_info = + ctx->callback_tracer->callback_data.at(info_type::callback_domain_idx); + callback_info.callback(record, &user_data, callback_info.data); + + // enter callback may update the external correlation id field + record.correlation_id.external = + ctx->correlation_tracer.external_correlator.get(thr_id); + } + } + + // record the start timestamp as close to the function call as possible + if(!buffered_contexts.empty()) + { + for(auto& itr : buffered_contexts) + { + itr.external_correlation = itr.ctx->correlation_tracer.external_correlator.get(thr_id); + } + + buffer_record.start_timestamp = common::timestamp_ns(); + } + + // decrement the reference count before invoking + corr_id->ref_count.fetch_sub(1); + + auto _ret = exec(info_type::get_table_func(), std::forward(args)...); + + // record the end timestamp as close to the function call as possible + if(!buffered_contexts.empty()) + { + buffer_record.end_timestamp = common::timestamp_ns(); + } + + if(!callback_contexts.empty()) + { + set_data_retval(tracer_data.retval, _ret); + + for(auto& itr : callback_contexts) + { + auto& ctx = itr.ctx; + auto& record = itr.record; + auto& user_data = itr.user_data; + + record.phase = ROCPROFILER_CALLBACK_PHASE_EXIT; + record.payload = static_cast(&tracer_data); + + auto& callback_info = + ctx->callback_tracer->callback_data.at(info_type::callback_domain_idx); + callback_info.callback(record, &user_data, callback_info.data); + } + } + + if(!buffered_contexts.empty()) + { + for(auto& itr : buffered_contexts) + { + assert(itr.ctx->buffered_tracer); + auto buffer_id = + itr.ctx->buffered_tracer->buffer_data.at(info_type::buffered_domain_idx); + auto buffer_v = buffer::get_buffer(buffer_id); + if(buffer_v && buffer_v->context_id == itr.ctx->context_idx && + buffer_v->buffer_id == buffer_id.handle) + { + // make copy of record + auto record_v = buffer_record; + // update the record with the correlation + record_v.correlation_id.external = itr.external_correlation; + + buffer_v->emplace( + ROCPROFILER_BUFFER_CATEGORY_TRACING, info_type::buffered_domain_idx, record_v); + } + } + } + + // decrement the reference count after usage in the callback/buffers + corr_id->ref_count.fetch_sub(1); + + context::pop_latest_correlation_id(corr_id); + + if constexpr(!std::is_same::value) + return _ret; + else + return 0; +} +} // namespace marker +} // namespace rocprofiler + +#define ROCPROFILER_LIB_ROCPROFILER_MARKER_MARKER_CPP_IMPL 1 + +// template specializations +#include "marker.def.cpp" + +namespace rocprofiler +{ +namespace marker +{ +namespace +{ +template +struct api_id_bounds; + +template <> +struct api_id_bounds +{ + static constexpr auto none = ROCPROFILER_MARKER_API_ID_NONE; + static constexpr auto last = ROCPROFILER_MARKER_API_ID_LAST; +}; + +template +const char* +name_by_id(const uint32_t id, std::index_sequence) +{ + if(OpIdx == id) return roctx_api_info::name; + + if constexpr(sizeof...(OpIdxTail) > 0) + return name_by_id(id, std::index_sequence{}); + else + return nullptr; +} + +template +uint32_t +id_by_name(const char* name, std::index_sequence) +{ + if(std::string_view{roctx_api_info::name} == std::string_view{name}) + return roctx_api_info::operation_idx; + + if constexpr(sizeof...(OpIdxTail) > 0) + return id_by_name(name, std::index_sequence{}); + else + return api_id_bounds::none; +} + +template +void +get_ids(std::vector& _id_list, std::index_sequence) +{ + auto _idx = roctx_api_info::operation_idx; + if(_idx < api_id_bounds::last) _id_list.emplace_back(_idx); + + if constexpr(sizeof...(OpIdxTail) > 0) + get_ids(_id_list, std::index_sequence{}); +} + +template +void +get_names(std::vector& _name_list, std::index_sequence) +{ + auto&& _name = roctx_api_info::name; + if(_name != nullptr && strnlen(_name, 1) > 0) _name_list.emplace_back(_name); + + if constexpr(sizeof...(OpIdxTail) > 0) + get_names(_name_list, std::index_sequence{}); +} + +template +void +iterate_args(const uint32_t id, + const rocprofiler_callback_tracing_marker_api_data_t& data, + rocprofiler_callback_tracing_operation_args_cb_t func, + void* user_data, + std::index_sequence) +{ + if(OpIdx == id) + { + using info_type = roctx_api_info; + auto&& arg_list = info_type::as_arg_list(data); + auto&& arg_addr = info_type::as_arg_addr(data); + for(size_t i = 0; i < std::min(arg_list.size(), arg_addr.size()); ++i) + { + auto ret = func(info_type::callback_domain_idx, // kind + id, // operation + i, // arg_number + arg_list.at(i).first.c_str(), // arg_name + arg_list.at(i).second.c_str(), // arg_value_str + arg_addr.at(i), // arg_value_addr + user_data); + if(ret != 0) break; + } + } + if constexpr(sizeof...(OpIdxTail) > 0) + iterate_args(id, data, func, user_data, std::index_sequence{}); +} + +bool +should_wrap_functor(rocprofiler_callback_tracing_kind_t _callback_domain, + rocprofiler_buffer_tracing_kind_t _buffered_domain, + int _operation) +{ + // we loop over all the *registered* contexts and see if any of them, at any point in time, + // might require callback or buffered API tracing + for(const auto& itr : context::get_registered_contexts()) + { + if(!itr) continue; + + // if there is a callback tracer enabled for the given domain and op, we need to wrap + if(itr->callback_tracer && itr->callback_tracer->domains(_callback_domain) && + itr->callback_tracer->domains(_callback_domain, _operation)) + return true; + + // if there is a buffered tracer enabled for the given domain and op, we need to wrap + if(itr->buffered_tracer && itr->buffered_tracer->domains(_buffered_domain) && + itr->buffered_tracer->domains(_buffered_domain, _operation)) + return true; + } + return false; +} + +template +void +copy_table(Tp* _orig, std::integral_constant) +{ + using table_type = typename roctx_table_lookup::type; + + if constexpr(std::is_same::value) + { + auto _info = roctx_api_info{}; + + LOG(INFO) << "copying table entry for " << _info.name; + + // make sure we don't access a field that doesn't exist in input table + if(_info.offset() >= _orig->size) return; + + // 1. get the sub-table containing the function pointer in original table + // 2. get reference to function pointer in sub-table in original table + auto& _table = _info.get_table(_orig); + auto& _func = _info.get_table_func(_table); + // 3. get the sub-table containing the function pointer in saved table + // 4. get reference to function pointer in sub-table in saved table + // 5. save the original function in the saved table + auto& _saved = _info.get_table(get_table()); + auto& _ofunc = _info.get_table_func(_saved); + _ofunc = _func; + } + + (void) _orig; +} + +template +void +update_table(Tp* _orig, std::integral_constant) +{ + using table_type = typename roctx_table_lookup::type; + + if constexpr(std::is_same::value) + { + auto _info = roctx_api_info{}; + + LOG(INFO) << "updating table entry for " << _info.name; + + // make sure we don't access a field that doesn't exist in input table + if(_info.offset() >= _orig->size) return; + + // check to see if there are any contexts which enable this operation in the ROCTX API + // domain + if(!should_wrap_functor( + _info.callback_domain_idx, _info.buffered_domain_idx, _info.operation_idx)) + return; + + // 1. get the sub-table containing the function pointer in original table + // 2. get reference to function pointer in sub-table in original table + // 3. update function pointer with wrapper + auto& _table = _info.get_table(_orig); + auto& _func = _info.get_table_func(_table); + _func = _info.get_functor(_func); + } + + (void) _orig; +} + +template +void +copy_table(Tp* _orig, std::index_sequence) +{ + copy_table(_orig, std::integral_constant{}); + if constexpr(sizeof...(OpIdxTail) > 0) + copy_table(_orig, std::index_sequence{}); +} + +template +void +update_table(Tp* _orig, std::index_sequence) +{ + update_table(_orig, std::integral_constant{}); + if constexpr(sizeof...(OpIdxTail) > 0) + update_table(_orig, std::index_sequence{}); +} +} // namespace + +// check out the assembly here... this compiles to a switch statement +template +const char* +name_by_id(uint32_t id) +{ + return name_by_id(id, std::make_index_sequence::last>{}); +} + +template +uint32_t +id_by_name(const char* name) +{ + return id_by_name(name, std::make_index_sequence::last>{}); +} + +template +std::vector +get_ids() +{ + constexpr auto last_api_id = api_id_bounds::last; + auto _data = std::vector{}; + _data.reserve(last_api_id); + get_ids(_data, std::make_index_sequence{}); + return _data; +} + +template +std::vector +get_names() +{ + constexpr auto last_api_id = api_id_bounds::last; + auto _data = std::vector{}; + _data.reserve(last_api_id); + get_names(_data, std::make_index_sequence{}); + return _data; +} + +void +iterate_args(uint32_t id, + const rocprofiler_callback_tracing_marker_api_data_t& data, + rocprofiler_callback_tracing_operation_args_cb_t callback, + void* user_data) +{ + if(callback) + iterate_args(id, + data, + callback, + user_data, + std::make_index_sequence{}); +} + +#define INSTANTIATE_MARKER_TABLE_FUNC(TABLE) \ + template const char* name_by_id(uint32_t); \ + template uint32_t id_by_name
(const char*); \ + template std::vector get_ids
(); \ + template std::vector get_names
(); + +INSTANTIATE_MARKER_TABLE_FUNC(ROCPROFILER_MARKER_API_TABLE_ID_RoctxApi) + +void +copy_table(roctx_api_table_t* _orig) +{ + if(_orig) + copy_table( + _orig, std::make_index_sequence{}); +} + +void +update_table(roctx_api_table_t* _orig) +{ + if(_orig) + update_table( + _orig, std::make_index_sequence{}); +} +} // namespace marker +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk/marker/marker.def.cpp b/source/lib/rocprofiler-sdk/marker/marker.def.cpp new file mode 100644 index 0000000000..aa589484f8 --- /dev/null +++ b/source/lib/rocprofiler-sdk/marker/marker.def.cpp @@ -0,0 +1,65 @@ +// 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/marker/defines.hpp" +#include "lib/rocprofiler-sdk/marker/marker.hpp" + +namespace rocprofiler +{ +namespace marker +{ +template <> +struct roctx_domain_info +{ + static constexpr auto callback_domain_idx = ROCPROFILER_CALLBACK_TRACING_MARKER_API; + static constexpr auto buffered_domain_idx = ROCPROFILER_BUFFER_TRACING_MARKER_API; + using args_type = rocprofiler_marker_api_args_t; + using retval_type = rocprofiler_marker_api_retval_t; + using callback_data_type = rocprofiler_callback_tracing_marker_api_data_t; +}; +} // namespace marker +} // namespace rocprofiler + +#if defined(ROCPROFILER_LIB_ROCPROFILER_MARKER_MARKER_CPP_IMPL) && \ + ROCPROFILER_LIB_ROCPROFILER_MARKER_MARKER_CPP_IMPL == 1 + +// clang-format off +MARKER_API_TABLE_LOOKUP_DEFINITION(ROCPROFILER_MARKER_API_TABLE_ID_RoctxApi, roctx_api_table_t) + +MARKER_API_INFO_DEFINITION_V(ROCPROFILER_MARKER_API_TABLE_ID_RoctxApi, ROCPROFILER_MARKER_API_ID_roctxMarkA, roctxMarkA, roctxMarkA_fn, message) +MARKER_API_INFO_DEFINITION_V(ROCPROFILER_MARKER_API_TABLE_ID_RoctxApi, ROCPROFILER_MARKER_API_ID_roctxRangePushA, roctxRangePushA, roctxRangePushA_fn, message) +MARKER_API_INFO_DEFINITION_0(ROCPROFILER_MARKER_API_TABLE_ID_RoctxApi, ROCPROFILER_MARKER_API_ID_roctxRangePop, roctxRangePop, roctxRangePop_fn) +MARKER_API_INFO_DEFINITION_V(ROCPROFILER_MARKER_API_TABLE_ID_RoctxApi, ROCPROFILER_MARKER_API_ID_roctxRangeStartA, roctxRangeStartA, roctxRangeStartA_fn, message) +MARKER_API_INFO_DEFINITION_V(ROCPROFILER_MARKER_API_TABLE_ID_RoctxApi, ROCPROFILER_MARKER_API_ID_roctxRangeStop, roctxRangeStop, roctxRangeStop_fn, id) +MARKER_API_INFO_DEFINITION_V(ROCPROFILER_MARKER_API_TABLE_ID_RoctxApi, ROCPROFILER_MARKER_API_ID_roctxProfilerPause, roctxProfilerPause, roctxProfilerPause_fn, tid) +MARKER_API_INFO_DEFINITION_V(ROCPROFILER_MARKER_API_TABLE_ID_RoctxApi, ROCPROFILER_MARKER_API_ID_roctxProfilerResume, roctxProfilerResume, roctxProfilerResume_fn, tid) +MARKER_API_INFO_DEFINITION_V(ROCPROFILER_MARKER_API_TABLE_ID_RoctxApi, ROCPROFILER_MARKER_API_ID_roctxNameOsThread, roctxNameOsThread, roctxNameOsThread_fn, name) +MARKER_API_INFO_DEFINITION_V(ROCPROFILER_MARKER_API_TABLE_ID_RoctxApi, ROCPROFILER_MARKER_API_ID_roctxNameHsaAgent, roctxNameHsaAgent, roctxNameHsaAgent_fn, name, agent) +MARKER_API_INFO_DEFINITION_V(ROCPROFILER_MARKER_API_TABLE_ID_RoctxApi, ROCPROFILER_MARKER_API_ID_roctxNameHipDevice, roctxNameHipDevice, roctxNameHipDevice_fn, name, device_id) +MARKER_API_INFO_DEFINITION_V(ROCPROFILER_MARKER_API_TABLE_ID_RoctxApi, ROCPROFILER_MARKER_API_ID_roctxNameHipStream, roctxNameHipStream, roctxNameHipStream_fn, name, stream) +MARKER_API_INFO_DEFINITION_V(ROCPROFILER_MARKER_API_TABLE_ID_RoctxApi, ROCPROFILER_MARKER_API_ID_roctxGetThreadId, roctxGetThreadId, roctxGetThreadId_fn, tid) +// clang-format on + +#else +# error \ + "Do not compile this file directly. It is included by lib/rocprofiler-sdk/marker/marker.cpp" +#endif diff --git a/source/lib/rocprofiler-sdk/marker/marker.hpp b/source/lib/rocprofiler-sdk/marker/marker.hpp new file mode 100644 index 0000000000..31d0bd01d3 --- /dev/null +++ b/source/lib/rocprofiler-sdk/marker/marker.hpp @@ -0,0 +1,91 @@ +// MIT License +// +// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +// THE SOFTWARE. + +#pragma once + +#include + +#include + +#include +#include + +namespace rocprofiler +{ +namespace marker +{ +using roctx_api_table_t = ::roctxApiTable_t; + +roctx_api_table_t& +get_table(); + +template +struct roctx_table_lookup; + +template +struct roctx_domain_info; + +template +struct roctx_api_info; + +template +struct roctx_api_impl : roctx_domain_info +{ + template + static auto set_data_args(DataArgsT&, Args... args); + + template + static auto exec(FuncT&&, Args&&... args); + + template + static auto functor(Args&&... args); +}; + +template +const char* +name_by_id(uint32_t id); + +template +uint32_t +id_by_name(const char* name); + +void +iterate_args(uint32_t id, + const rocprofiler_callback_tracing_marker_api_data_t& data, + rocprofiler_callback_tracing_operation_args_cb_t callback, + void* user_data); + +template +std::vector +get_names(); + +template +std::vector +get_ids(); + +void +copy_table(roctx_api_table_t* _orig); + +void +update_table(roctx_api_table_t* _orig); +} // namespace marker +} // namespace rocprofiler diff --git a/source/lib/rocprofiler-sdk/marker/utils.hpp b/source/lib/rocprofiler-sdk/marker/utils.hpp new file mode 100644 index 0000000000..224f718ab9 --- /dev/null +++ b/source/lib/rocprofiler-sdk/marker/utils.hpp @@ -0,0 +1,124 @@ +// MIT License +// +// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +// THE SOFTWARE. + +#pragma once + +#include +#include + +#include "lib/common/mpl.hpp" + +#include "fmt/core.h" +#include "fmt/ranges.h" + +#include +#include +#include +#include +#include +#include + +namespace rocprofiler +{ +namespace marker +{ +namespace utils +{ +template +auto +stringize_impl(const Tp& _v) +{ + using nonpointer_type = typename std::remove_pointer_t; + + if constexpr(common::mpl::is_pair::value) + { + return std::make_pair(stringize_impl(_v.first), stringize_impl(_v.second)); + } + else if constexpr(std::is_constructible::value) + { + auto _ss = std::stringstream{}; + _ss << _v; + return _ss.str(); + } + else if constexpr(fmt::is_formattable::value && !std::is_pointer::value) + { + return fmt::format("{}", _v); + } + else if constexpr(std::is_pointer::value && !std::is_pointer::value && + common::mpl::is_type_complete_v && + !std::is_void::value) + { + if(_v) + { + return stringize_impl(*_v); + } + else + { + auto _ss = std::stringstream{}; + _ss << _v; + return _ss.str(); + } + } + else + { + auto _ss = std::stringstream{}; + _ss << _v; + return _ss.str(); + } +} + +template +auto +stringize(Args... args) +{ + return std::vector>{stringize_impl(args)...}; +} + +template +struct handle_formatter +{ + template + constexpr auto parse(ParseContext& ctx) + { + return ctx.begin(); + } + + template + auto format(const Tp& v, Ctx& ctx) const + { + return fmt::format_to(ctx.out(), "handle={}", v.handle); + } +}; + +template +struct handle_formatter : handle_formatter +{}; +} // namespace utils +} // namespace marker +} // namespace rocprofiler + +namespace fmt +{ +template <> +struct formatter : rocprofiler::marker::utils::handle_formatter +{}; +} // namespace fmt diff --git a/source/lib/rocprofiler-sdk/registration.cpp b/source/lib/rocprofiler-sdk/registration.cpp index c0de2dc14e..61e61887b6 100644 --- a/source/lib/rocprofiler-sdk/registration.cpp +++ b/source/lib/rocprofiler-sdk/registration.cpp @@ -33,6 +33,7 @@ #include "lib/rocprofiler-sdk/hsa/queue_controller.hpp" #include "lib/rocprofiler-sdk/intercept_table.hpp" #include "lib/rocprofiler-sdk/internal_threading.hpp" +#include "lib/rocprofiler-sdk/marker/marker.hpp" #include #include @@ -156,7 +157,7 @@ struct client_library rocprofiler_client_id_t mutable_client_id = {}; }; -using client_library_vec_t = std::vector; +using client_library_vec_t = std::vector>; client_library_vec_t find_clients() @@ -245,17 +246,17 @@ find_clients() // skip the configure function that was forced if(_sym == get_forced_configure()) { - data.front().name = itr; - data.front().dlhandle = handle; - data.front().internal_client_id.name = "(forced)"; + data.front()->name = itr; + data.front()->dlhandle = handle; + data.front()->internal_client_id.name = "(forced)"; continue; } if(_sym == &rocprofiler_configure && data.size() == 1) { - data.front().name = itr; - data.front().dlhandle = handle; - data.front().internal_client_id.name = "default"; + data.front()->name = itr; + data.front()->dlhandle = handle; + data.front()->internal_client_id.name = "default"; } else { @@ -267,7 +268,7 @@ find_clients() nullptr, rocprofiler_client_id_t{nullptr, _prio}, rocprofiler_client_id_t{nullptr, _prio}}); - entry.internal_client_id.name = entry.name.c_str(); + entry->internal_client_id.name = entry->name.c_str(); } } @@ -280,7 +281,7 @@ find_clients() for(const auto& ditr : data) { - if(ditr.dlhandle && ditr.dlhandle == handle) + if(ditr->dlhandle && ditr->dlhandle == handle) { handle = nullptr; break; @@ -315,7 +316,7 @@ get_clients() return _v; } -using mutex_t = std::recursive_mutex; +using mutex_t = std::mutex; using scoped_lock_t = std::unique_lock; mutex_t& @@ -330,9 +331,7 @@ invoke_client_configures() { if(get_init_status() > 0) return false; - auto _lk = scoped_lock_t{get_registration_mutex(), std::defer_lock}; - if(_lk.owns_lock()) return false; - _lk.lock(); + auto _lk = scoped_lock_t{get_registration_mutex()}; LOG(ERROR) << __FUNCTION__; @@ -340,20 +339,22 @@ invoke_client_configures() for(auto& itr : *get_clients()) { - if(!itr.configure_func) + if(!itr) continue; + + if(!itr->configure_func) { LOG(ERROR) << "rocprofiler::registration::invoke_client_configures() attempted to " "invoke configure function from " - << itr.name << " that had no configuration function"; + << itr->name << " that had no configuration function"; continue; } - if(get_invoked_configures().find(itr.configure_func) != get_invoked_configures().end()) + if(get_invoked_configures().find(itr->configure_func) != get_invoked_configures().end()) { LOG(ERROR) << "rocprofiler::registration::invoke_client_configures() attempted to " "invoke configure function from " - << itr.name << " (addr=" - << fmt::format("{:#018x}", reinterpret_cast(itr.configure_func)) + << itr->name << " (addr=" + << fmt::format("{:#018x}", reinterpret_cast(itr->configure_func)) << ") more than once"; continue; } @@ -361,27 +362,27 @@ invoke_client_configures() { LOG(INFO) << "rocprofiler::registration::invoke_client_configures() invoking configure " "function from " - << itr.name << " (addr=" - << fmt::format("{:#018x}", reinterpret_cast(itr.configure_func)) + << itr->name << " (addr=" + << fmt::format("{:#018x}", reinterpret_cast(itr->configure_func)) << ")"; } - auto* _result = itr.configure_func(ROCPROFILER_VERSION, - ROCPROFILER_VERSION_STRING, - itr.internal_client_id.handle - get_client_offset(), - &itr.mutable_client_id); + auto* _result = itr->configure_func(ROCPROFILER_VERSION, + ROCPROFILER_VERSION_STRING, + itr->internal_client_id.handle - get_client_offset(), + &itr->mutable_client_id); if(_result) { - itr.configure_result = new rocprofiler_tool_configure_result_t{*_result}; + itr->configure_result = new rocprofiler_tool_configure_result_t{*_result}; } else { - context::deactivate_client_contexts(itr.internal_client_id); - context::deregister_client_contexts(itr.internal_client_id); + context::deactivate_client_contexts(itr->internal_client_id); + context::deregister_client_contexts(itr->internal_client_id); } - get_invoked_configures().emplace(itr.configure_func); + get_invoked_configures().emplace(itr->configure_func); } return true; @@ -392,9 +393,7 @@ invoke_client_initializers() { if(get_init_status() > 0) return false; - auto _lk = scoped_lock_t{get_registration_mutex(), std::defer_lock}; - if(_lk.owns_lock()) return false; - _lk.lock(); + auto _lk = scoped_lock_t{get_registration_mutex()}; LOG(ERROR) << __FUNCTION__; @@ -402,14 +401,14 @@ invoke_client_initializers() for(auto& itr : *get_clients()) { - if(itr.configure_result && itr.configure_result->initialize) + if(itr && itr->configure_result && itr->configure_result->initialize) { - context::push_client(itr.internal_client_id.handle); - itr.configure_result->initialize(&invoke_client_finalizer, - itr.configure_result->tool_data); - context::pop_client(itr.internal_client_id.handle); + context::push_client(itr->internal_client_id.handle); + itr->configure_result->initialize(&invoke_client_finalizer, + itr->configure_result->tool_data); + context::pop_client(itr->internal_client_id.handle); // set to nullptr so initialize only gets called once - itr.configure_result->initialize = nullptr; + itr->configure_result->initialize = nullptr; } } @@ -424,23 +423,12 @@ invoke_client_finalizers() if(get_init_status() < 1 || get_fini_status() > 0) return false; - auto _lk = scoped_lock_t{get_registration_mutex(), std::defer_lock}; - if(_lk.owns_lock()) return false; - _lk.lock(); - - LOG(ERROR) << __FUNCTION__; - - if(!get_clients()) return false; - - for(auto& itr : *get_clients()) + if(get_clients()) { - if(itr.configure_result && itr.configure_result->finalize) + for(auto& itr : *get_clients()) { - itr.configure_result->finalize(itr.configure_result->tool_data); - // set to nullptr so finalize only gets called once - itr.configure_result->finalize = nullptr; + if(itr) invoke_client_finalizer(itr->internal_client_id); } - context::deactivate_client_contexts(itr.internal_client_id); } return true; @@ -449,29 +437,30 @@ invoke_client_finalizers() void invoke_client_finalizer(rocprofiler_client_id_t client_id) { - auto _lk = scoped_lock_t{get_registration_mutex(), std::defer_lock}; - if(_lk.owns_lock()) return; - _lk.lock(); + LOG(ERROR) << __FUNCTION__ << "(client_id=" << client_id.handle << ")"; - LOG(ERROR) << __FUNCTION__; + auto _lk = scoped_lock_t{get_registration_mutex()}; if(!get_clients()) return; for(auto& itr : *get_clients()) { - if(itr.internal_client_id.handle == client_id.handle && - itr.mutable_client_id.handle == client_id.handle) + if(itr && itr->internal_client_id.handle == client_id.handle && + itr->mutable_client_id.handle == client_id.handle) { - if(itr.configure_result && itr.configure_result->finalize) + if(itr->configure_result && itr->configure_result->finalize) { + // set to nullptr so finalize only gets called once + rocprofiler_tool_finalize_t _finalize_func = nullptr; + std::swap(_finalize_func, itr->configure_result->finalize); + auto _fini_status = get_fini_status(); if(_fini_status == 0) set_fini_status(-1); - itr.configure_result->finalize(itr.configure_result->tool_data); + _finalize_func(itr->configure_result->tool_data); if(_fini_status == 0) set_fini_status(_fini_status); - // set to nullptr so finalize only gets called once - itr.configure_result->finalize = nullptr; } - context::deactivate_client_contexts(itr.internal_client_id); + context::deactivate_client_contexts(itr->internal_client_id); + itr.reset(); } } } @@ -684,7 +673,21 @@ rocprofiler_set_api_table(const char* name, LOG_IF(ERROR, num_tables > 1) << " rocprofiler expected ROCTX library to pass 1 API table, not " << num_tables; - return ROCPROFILER_STATUS_ERROR_NOT_IMPLEMENTED; + auto* roctx_api_table = static_cast(*tables); + + // any internal modifications to the roctxApiTable_t need to be done before we make + // the copy or else those modifications will be lost when ROCTx tracing is enabled because + // the ROCTx tracing invokes the function pointers from the copy below + rocprofiler::marker::copy_table(roctx_api_table); + + // install rocprofiler API wrappers + rocprofiler::marker::update_table(roctx_api_table); + + rocprofiler::intercept_table::notify_runtime_api_registration( + ROCPROFILER_MARKER_LIBRARY, + lib_version, + lib_instance, + std::make_tuple(roctx_api_table)); } else { diff --git a/source/lib/rocprofiler-sdk/tests/CMakeLists.txt b/source/lib/rocprofiler-sdk/tests/CMakeLists.txt index 6377245119..1cf57a3d2f 100644 --- a/source/lib/rocprofiler-sdk/tests/CMakeLists.txt +++ b/source/lib/rocprofiler-sdk/tests/CMakeLists.txt @@ -36,15 +36,20 @@ set_tests_properties(${lib_TESTS} PROPERTIES TIMEOUT 45 LABELS "unittests") # -------------------------------------------------------------------------------------- # set(rocprofiler_shared_lib_sources external_correlation.cpp intercept_table.cpp - registration.cpp status.cpp) + registration.cpp roctx.cpp status.cpp) add_executable(rocprofiler-lib-tests-shared) target_sources(rocprofiler-lib-tests-shared PRIVATE ${rocprofiler_shared_lib_sources}) target_link_libraries( rocprofiler-lib-tests-shared PRIVATE rocprofiler::rocprofiler-shared-library - rocprofiler::rocprofiler-common-library rocprofiler::rocprofiler-hsa-runtime - GTest::gtest GTest::gtest_main) + rocprofiler::rocprofiler-common-library + rocprofiler::rocprofiler-hsa-runtime + rocprofiler::rocprofiler-sdk-roctx-shared-library + GTest::gtest + GTest::gtest_main) +set_target_properties(rocprofiler-lib-tests-shared PROPERTIES BUILD_RPATH + "\$ORIGIN/../lib") gtest_add_tests( TARGET rocprofiler-lib-tests-shared @@ -52,7 +57,10 @@ gtest_add_tests( TEST_LIST shared_lib_TESTS WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}) +set(rocprofiler-lib-tests-env + "HSA_TOOLS_LIB=$" + "LD_LIBRARY_PATH=${PROJECT_BINARY_DIR}/${CMAKE_INSTALL_LIBDIR}:$ENV{LD_LIBRARY_PATH}") + set_tests_properties( - ${shared_lib_TESTS} - PROPERTIES TIMEOUT 360 LABELS "unittests" ENVIRONMENT - "HSA_TOOLS_LIB=$") + ${shared_lib_TESTS} PROPERTIES TIMEOUT 360 LABELS "unittests" ENVIRONMENT + "${rocprofiler-lib-tests-env}") diff --git a/source/lib/rocprofiler-sdk/tests/common.hpp b/source/lib/rocprofiler-sdk/tests/common.hpp index ae4c44f4ea..9d2e2864c0 100644 --- a/source/lib/rocprofiler-sdk/tests/common.hpp +++ b/source/lib/rocprofiler-sdk/tests/common.hpp @@ -25,6 +25,7 @@ #include #include "lib/common/defines.hpp" +#include "rocprofiler-sdk/fwd.h" #include @@ -76,7 +77,7 @@ inline auto get_callback_tracing_names() { static const auto supported_kinds = std::unordered_set{ - ROCPROFILER_CALLBACK_TRACING_HSA_API}; + ROCPROFILER_CALLBACK_TRACING_HSA_API, ROCPROFILER_CALLBACK_TRACING_MARKER_API}; auto cb_name_info = callback_name_info{}; // @@ -139,8 +140,8 @@ struct buffer_name_info inline buffer_name_info get_buffer_tracing_names() { - static const auto supported_kinds = - std::unordered_set{ROCPROFILER_BUFFER_TRACING_HSA_API}; + static const auto supported_kinds = std::unordered_set{ + ROCPROFILER_BUFFER_TRACING_HSA_API, ROCPROFILER_BUFFER_TRACING_MARKER_API}; auto cb_name_info = buffer_name_info{}; // diff --git a/source/lib/rocprofiler-sdk/tests/roctx.cpp b/source/lib/rocprofiler-sdk/tests/roctx.cpp new file mode 100644 index 0000000000..468953be0e --- /dev/null +++ b/source/lib/rocprofiler-sdk/tests/roctx.cpp @@ -0,0 +1,542 @@ +// 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 +#include +#include +#include +#include +#include + +#include "lib/common/defines.hpp" +#include "lib/common/environment.hpp" +#include "lib/common/filesystem.hpp" +#include "lib/common/units.hpp" +#include "lib/common/utility.hpp" +#include "lib/rocprofiler-sdk/tests/common.hpp" +#include "rocprofiler-sdk-roctx/api_trace.h" +#include "rocprofiler-sdk-roctx/types.h" +#include "rocprofiler-sdk/callback_tracing.h" +#include "rocprofiler-sdk/context.h" +#include "rocprofiler-sdk/marker/api_id.h" + +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace +{ +#define ROCTX_CALL(ARG, MSG) \ + { \ + auto _status = (ARG); \ + EXPECT_EQ(_status, 0) << MSG << " :: " << #ARG; \ + } + +void +run_roctx_functions() +{ + auto roctx_tid = roctx_thread_id_t{}; + auto rocp_tid = rocprofiler_thread_id_t{}; + auto hsa_agent = hsa_agent_t{.handle = roctx_tid}; + auto* hip_stream = hipStream_t{}; + + roctxGetThreadId(&roctx_tid); + rocprofiler_get_thread_id(&rocp_tid); + + EXPECT_EQ(roctx_tid, rocp_tid); + roctxNameOsThread(std::to_string(roctx_tid).c_str()); + roctxNameHsaAgent("hsa_agent", &hsa_agent); + roctxNameHipDevice("hip_device", 0); + roctxNameHipStream("hip_device", hip_stream); + + roctxProfilerPause(roctx_tid); + roctxMark(__FUNCTION__); + roctxProfilerResume(roctx_tid); + + ROCTX_CALL(roctxRangePush(__FUNCTION__), "push failed"); + ROCTX_CALL(roctxRangePop(), "pop failed"); + + auto range_id = roctxRangeStart(__FUNCTION__); + EXPECT_GE(range_id, 0); + roctxRangeStop(range_id); +} + +void +tool_tracing_ctrl_callback(rocprofiler_callback_tracing_record_t record, + rocprofiler_user_data_t*, + void* client_data) +{ + auto* cb_data = static_cast(client_data); + + if(record.phase == ROCPROFILER_CALLBACK_PHASE_ENTER && + record.kind == ROCPROFILER_CALLBACK_TRACING_MARKER_API && + record.operation == ROCPROFILER_MARKER_API_ID_roctxProfilerPause) + { + ROCPROFILER_CALL(rocprofiler_stop_context(cb_data->client_ctx), "pausing client context"); + } + else if(record.phase == ROCPROFILER_CALLBACK_PHASE_EXIT && + record.kind == ROCPROFILER_CALLBACK_TRACING_MARKER_API && + record.operation == ROCPROFILER_MARKER_API_ID_roctxProfilerResume) + { + ROCPROFILER_CALL(rocprofiler_start_context(cb_data->client_ctx), "resuming client context"); + } +} + +void +tool_tracing_callback(rocprofiler_callback_tracing_record_t record, + rocprofiler_user_data_t* user_data, + void* client_data) +{ + auto* cb_data = static_cast(client_data); + auto get_timestamp = []() { + return std::chrono::steady_clock::now().time_since_epoch().count(); + }; + + if(record.phase == ROCPROFILER_CALLBACK_PHASE_ENTER && cb_data->current_depth == 0) + { + user_data->value = get_timestamp(); + } + + static auto name_map = get_callback_tracing_names(); + + EXPECT_EQ(name_map.kind_names.size(), ROCPROFILER_CALLBACK_TRACING_LAST); + EXPECT_EQ(name_map.operation_names.at(ROCPROFILER_CALLBACK_TRACING_MARKER_API).size(), + ROCPROFILER_MARKER_API_ID_LAST); + + std::cout << "[" << __FILE__ << ":" << __LINE__ << "] " + << name_map.operation_names[record.kind][record.operation] << "\n" + << std::flush; + + cb_data->client_callback_count++; + if(record.phase == ROCPROFILER_CALLBACK_PHASE_ENTER) + { + cb_data->current_depth++; + } + else if(record.phase == ROCPROFILER_CALLBACK_PHASE_EXIT) + { + cb_data->max_depth = std::max(cb_data->current_depth, cb_data->max_depth); + cb_data->current_depth--; + } + else + { + GTEST_FAIL() << "unsupported callback tracing phase " << record.phase; + } + + struct info_data + { + uint64_t num_args = 0; + std::stringstream arg_ss; + } info_data_v; + + auto info_data_cb = [](rocprofiler_callback_tracing_kind_t, + uint32_t, + uint32_t arg_num, + const char* arg_name, + const char* arg_value_str, + const void* const arg_value_addr, + void* data) -> int { + auto& info = *static_cast(data); + info.arg_ss << ((arg_num == 0) ? "(" : ", "); + info.arg_ss << arg_num << ": " << arg_name << "=" << arg_value_str; + EXPECT_NE(arg_name, nullptr); + EXPECT_NE(arg_value_str, nullptr); + EXPECT_NE(arg_value_addr, nullptr); + EXPECT_EQ(arg_num, info.num_args); + info.num_args++; + return 0; + }; + + ROCPROFILER_CALL(rocprofiler_iterate_callback_tracing_kind_operation_args( + record, info_data_cb, static_cast(&info_data_v)), + "Failure iterating trace operation args"); + if(record.kind == ROCPROFILER_CALLBACK_TRACING_MARKER_API) + { + if(record.operation != ROCPROFILER_MARKER_API_ID_roctxRangePop) + { + EXPECT_GT(info_data_v.num_args, 0) + << name_map.operation_names[record.kind][record.operation] + << info_data_v.arg_ss.str(); + } + } + + if(record.phase == ROCPROFILER_CALLBACK_PHASE_EXIT && cb_data->current_depth == 0) + { + cb_data->client_elapsed += (get_timestamp() - user_data->value); + } +} + +void +tool_tracing_buffered(rocprofiler_context_id_t context, + rocprofiler_buffer_id_t buffer_id, + rocprofiler_record_header_t** headers, + size_t num_headers, + void* buffer_data, + uint64_t drop_count) +{ + auto* cb_data = static_cast(buffer_data); + + static auto name_map = get_buffer_tracing_names(); + + EXPECT_EQ(name_map.kind_names.size(), ROCPROFILER_BUFFER_TRACING_LAST); + EXPECT_EQ(name_map.operation_names.at(ROCPROFILER_BUFFER_TRACING_MARKER_API).size(), + ROCPROFILER_MARKER_API_ID_LAST); + + auto v_records = std::vector{}; + v_records.reserve(num_headers); + + for(size_t i = 0; i < num_headers; ++i) + { + auto* header = headers[i]; + + ASSERT_TRUE(header != nullptr); + auto hash = rocprofiler_record_header_compute_hash(header->category, header->kind); + EXPECT_EQ(header->hash, hash); + EXPECT_TRUE(header->category == ROCPROFILER_BUFFER_CATEGORY_TRACING && + header->kind == ROCPROFILER_BUFFER_TRACING_MARKER_API); + + v_records.emplace_back( + static_cast(header->payload)); + } + + std::sort(v_records.begin(), v_records.end(), [](auto lhs, auto rhs) { + return (lhs->start_timestamp == rhs->start_timestamp) + ? (lhs->end_timestamp < rhs->end_timestamp) + : (lhs->start_timestamp < rhs->start_timestamp); + }); + + for(auto* record : v_records) + { + auto info = std::stringstream{}; + info << "tid=" << record->thread_id << ", context=" << context.handle + << ", buffer_id=" << buffer_id.handle << ", cid=" << record->correlation_id.internal + << ", kind=" << name_map.kind_names.at(record->kind) << "(" << record->kind + << "), operation=" << name_map.operation_names.at(record->kind).at(record->operation) + << "(" << record->operation << "), drop_count=" << drop_count + << ", start=" << record->start_timestamp << ", stop=" << record->end_timestamp; + + static int64_t last_corr_id = -1; + auto corr_id = static_cast(record->correlation_id.internal); + + std::cout << info.str() << "\n" << std::flush; + EXPECT_GE(context.handle, 0) << info.str(); + EXPECT_GT(record->thread_id, 0) << info.str(); + EXPECT_GT(record->kind, 0) << info.str(); + EXPECT_GT(corr_id, last_corr_id) << info.str(); + EXPECT_GT(record->start_timestamp, 0) << info.str(); + EXPECT_GT(record->end_timestamp, 0) << info.str(); + EXPECT_LE(record->start_timestamp, record->end_timestamp) << info.str(); + + cb_data->client_callback_count++; + last_corr_id = corr_id; + } +} + +void +thread_precreate(rocprofiler_runtime_library_t /*lib*/, void* tool_data) +{ + auto* cb_data = static_cast(tool_data); + cb_data->client_workflow_count++; +} + +void +thread_postcreate(rocprofiler_runtime_library_t /*lib*/, void* tool_data) +{ + auto* cb_data = static_cast(tool_data); + cb_data->client_workflow_count++; +} +} // namespace + +TEST(rocprofiler_lib, roctx_callback_tracing) +{ + using init_func_t = int (*)(rocprofiler_client_finalize_t, void*); + using fini_func_t = void (*)(void*); + + auto cmd_line = rocprofiler::common::read_command_line(getpid()); + + ASSERT_FALSE(cmd_line.empty()); + + static init_func_t tool_init = [](rocprofiler_client_finalize_t fini_func, + void* client_data) -> int { + auto* cb_data = static_cast(client_data); + + cb_data->client_workflow_count++; + cb_data->client_fini_func = fini_func; + + auto pause_resume_ctx = rocprofiler_context_id_t{}; + ROCPROFILER_CALL(rocprofiler_create_context(&pause_resume_ctx), "failed to create context"); + + ROCPROFILER_CALL(rocprofiler_create_context(&cb_data->client_ctx), + "failed to create context"); + + auto operations = std::vector{}; + rocprofiler_iterate_callback_tracing_kind_operations( + ROCPROFILER_CALLBACK_TRACING_MARKER_API, + [](rocprofiler_callback_tracing_kind_t, uint32_t operation, void* data) { + auto* _ops = static_cast*>(data); + if(operation != ROCPROFILER_MARKER_API_ID_roctxProfilerPause && + operation != ROCPROFILER_MARKER_API_ID_roctxProfilerResume) + _ops->emplace_back(operation); + return 0; + }, + &operations); + + ROCPROFILER_CALL( + rocprofiler_configure_callback_tracing_service(cb_data->client_ctx, + ROCPROFILER_CALLBACK_TRACING_MARKER_API, + operations.data(), + operations.size(), + tool_tracing_callback, + client_data), + "callback tracing service failed to configure"); + + auto pause_resume_ops = + std::array{ROCPROFILER_MARKER_API_ID_roctxProfilerPause, + ROCPROFILER_MARKER_API_ID_roctxProfilerResume}; + + ROCPROFILER_CALL( + rocprofiler_configure_callback_tracing_service(pause_resume_ctx, + ROCPROFILER_CALLBACK_TRACING_MARKER_API, + pause_resume_ops.data(), + pause_resume_ops.size(), + tool_tracing_ctrl_callback, + client_data), + "callback tracing service failed to configure"); + + int valid_ctx = 0; + ROCPROFILER_CALL(rocprofiler_context_is_valid(cb_data->client_ctx, &valid_ctx), + "failure checking context validity"); + + EXPECT_EQ(valid_ctx, 1); + + ROCPROFILER_CALL(rocprofiler_start_context(pause_resume_ctx), + "rocprofiler context start failed"); + + ROCPROFILER_CALL(rocprofiler_start_context(cb_data->client_ctx), + "rocprofiler context start failed"); + + // no errors + return 0; + }; + + static fini_func_t tool_fini = [](void* client_data) -> void { + auto* cb_data = static_cast(client_data); + ROCPROFILER_CALL(rocprofiler_stop_context(cb_data->client_ctx), + "rocprofiler context stop failed"); + + static_cast(client_data)->client_workflow_count++; + }; + + static auto cb_data = callback_data{}; + + static auto cfg_result = + rocprofiler_tool_configure_result_t{sizeof(rocprofiler_tool_configure_result_t), + tool_init, + tool_fini, + static_cast(&cb_data)}; + + static rocprofiler_configure_func_t rocp_init = + [](uint32_t version, + const char* runtime_version, + uint32_t prio, + rocprofiler_client_id_t* client_id) -> rocprofiler_tool_configure_result_t* { + auto expected_version = ROCPROFILER_VERSION; + EXPECT_EQ(expected_version, version); + EXPECT_EQ(std::string_view{runtime_version}, std::string_view{ROCPROFILER_VERSION_STRING}); + EXPECT_EQ(prio, 0); + EXPECT_EQ(client_id->name, nullptr); + cb_data.client_id = client_id; + cb_data.client_id->name = ::testing::UnitTest::GetInstance()->current_test_info()->name(); + return &cfg_result; + }; + + EXPECT_EQ(rocprofiler_force_configure(rocp_init), ROCPROFILER_STATUS_SUCCESS); + + run_roctx_functions(); + + ASSERT_NE(cb_data.client_id, nullptr); + ASSERT_NE(cb_data.client_fini_func, nullptr); + + cb_data.client_fini_func(*cb_data.client_id); + + // expected callback count is two for hsa_iterate_agents and two callbacks for + // hsa_agent_get_info for each agent. + constexpr uint64_t expected_cb_count = 18; + + EXPECT_EQ(cb_data.client_workflow_count, 2); + EXPECT_EQ(cb_data.client_callback_count, expected_cb_count); + EXPECT_EQ(cb_data.current_depth, 0); + EXPECT_EQ(cb_data.max_depth, 1); +} + +TEST(rocprofiler_lib, roctx_buffered_tracing) +{ + using init_func_t = int (*)(rocprofiler_client_finalize_t, void*); + using fini_func_t = void (*)(void*); + + auto cmd_line = rocprofiler::common::read_command_line(getpid()); + ASSERT_FALSE(cmd_line.empty()); + + static init_func_t tool_init = [](rocprofiler_client_finalize_t fini_func, + void* client_data) -> int { + auto* cb_data = static_cast(client_data); + + cb_data->client_workflow_count++; + cb_data->client_fini_func = fini_func; + + auto pause_resume_ctx = rocprofiler_context_id_t{}; + auto pause_resume_ops = + std::array{ROCPROFILER_MARKER_API_ID_roctxProfilerPause, + ROCPROFILER_MARKER_API_ID_roctxProfilerResume}; + + ROCPROFILER_CALL(rocprofiler_create_context(&pause_resume_ctx), "failed to create context"); + ROCPROFILER_CALL( + rocprofiler_configure_callback_tracing_service(pause_resume_ctx, + ROCPROFILER_CALLBACK_TRACING_MARKER_API, + pause_resume_ops.data(), + pause_resume_ops.size(), + tool_tracing_ctrl_callback, + client_data), + "callback tracing service failed to configure"); + ROCPROFILER_CALL(rocprofiler_start_context(pause_resume_ctx), + "rocprofiler context start failed"); + + ROCPROFILER_CALL(rocprofiler_create_context(&cb_data->client_ctx), + "failed to create context"); + + ROCPROFILER_CALL(rocprofiler_create_buffer(cb_data->client_ctx, + 4096, + 2048, + ROCPROFILER_BUFFER_POLICY_LOSSLESS, + tool_tracing_buffered, + client_data, + &cb_data->client_buffer), + "buffer creation failed"); + + auto operations = std::vector{}; + rocprofiler_iterate_buffer_tracing_kind_operations( + ROCPROFILER_BUFFER_TRACING_MARKER_API, + [](rocprofiler_buffer_tracing_kind_t, uint32_t operation, void* data) { + auto* _ops = static_cast*>(data); + if(operation != ROCPROFILER_MARKER_API_ID_roctxProfilerPause && + operation != ROCPROFILER_MARKER_API_ID_roctxProfilerResume) + _ops->emplace_back(operation); + return 0; + }, + &operations); + + ROCPROFILER_CALL( + rocprofiler_configure_buffer_tracing_service(cb_data->client_ctx, + ROCPROFILER_BUFFER_TRACING_MARKER_API, + operations.data(), + operations.size(), + cb_data->client_buffer), + "buffer tracing service failed to configure"); + + ROCPROFILER_CALL(rocprofiler_create_callback_thread(&cb_data->client_thread), + "failure creating callback thread"); + + ROCPROFILER_CALL( + rocprofiler_assign_callback_thread(cb_data->client_buffer, cb_data->client_thread), + "failed to assign thread for buffer"); + + int valid_ctx = 0; + ROCPROFILER_CALL(rocprofiler_context_is_valid(cb_data->client_ctx, &valid_ctx), + "failure checking context validity"); + + EXPECT_EQ(valid_ctx, 1); + + ROCPROFILER_CALL(rocprofiler_start_context(cb_data->client_ctx), + "rocprofiler context start failed"); + + // no errors + return 0; + }; + + static fini_func_t tool_fini = [](void* client_data) -> void { + auto* cb_data = static_cast(client_data); + ROCPROFILER_CALL(rocprofiler_flush_buffer(cb_data->client_buffer), + "rocprofiler context stop failed"); + ROCPROFILER_CALL(rocprofiler_stop_context(cb_data->client_ctx), + "rocprofiler context stop failed"); + + static_cast(client_data)->client_workflow_count++; + }; + + static auto cb_data = callback_data{}; + + static auto cfg_result = + rocprofiler_tool_configure_result_t{sizeof(rocprofiler_tool_configure_result_t), + tool_init, + tool_fini, + static_cast(&cb_data)}; + + static rocprofiler_configure_func_t rocp_init = + [](uint32_t version, + const char* runtime_version, + uint32_t prio, + rocprofiler_client_id_t* client_id) -> rocprofiler_tool_configure_result_t* { + auto expected_version = ROCPROFILER_VERSION; + EXPECT_EQ(expected_version, version); + EXPECT_EQ(std::string_view{runtime_version}, std::string_view{ROCPROFILER_VERSION_STRING}); + EXPECT_EQ(prio, 0); + EXPECT_EQ(client_id->name, nullptr); + cb_data.client_id = client_id; + cb_data.client_id->name = ::testing::UnitTest::GetInstance()->current_test_info()->name(); + + ROCPROFILER_CALL(rocprofiler_at_internal_thread_create(thread_precreate, + thread_postcreate, + ROCPROFILER_LIBRARY, + static_cast(&cb_data)), + "failed to register for thread creation notifications"); + + return &cfg_result; + }; + + EXPECT_EQ(rocprofiler_force_configure(rocp_init), ROCPROFILER_STATUS_SUCCESS); + + run_roctx_functions(); + + ASSERT_NE(cb_data.client_id, nullptr); + ASSERT_NE(cb_data.client_fini_func, nullptr); + + cb_data.client_fini_func(*cb_data.client_id); + + constexpr uint64_t expected_cb_count = 9; + + EXPECT_EQ(cb_data.client_workflow_count, 6); + EXPECT_EQ(cb_data.client_callback_count, expected_cb_count); + EXPECT_EQ(cb_data.current_depth, 0); + EXPECT_EQ(cb_data.max_depth, 0); +} diff --git a/tests/apps/CMakeLists.txt b/tests/apps/CMakeLists.txt index 5d2061ab75..c5ef92a52c 100644 --- a/tests/apps/CMakeLists.txt +++ b/tests/apps/CMakeLists.txt @@ -5,6 +5,10 @@ cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR) project(rocprofiler-test-apps LANGUAGES C CXX) +set(CMAKE_BUILD_RPATH + "\$ORIGIN:\$ORIGIN/../lib:$" + ) + # applications used by integration tests add_subdirectory(reproducible-runtime) add_subdirectory(transpose) diff --git a/tests/apps/reproducible-runtime/CMakeLists.txt b/tests/apps/reproducible-runtime/CMakeLists.txt index 0a58bd9f38..23ee5583a4 100644 --- a/tests/apps/reproducible-runtime/CMakeLists.txt +++ b/tests/apps/reproducible-runtime/CMakeLists.txt @@ -43,6 +43,10 @@ target_compile_options(reproducible-runtime PRIVATE -W -Wall -Wextra -Wpedantic find_package(Threads REQUIRED) target_link_libraries(reproducible-runtime PRIVATE Threads::Threads) +find_package(rocprofiler-sdk-roctx REQUIRED) +target_link_libraries(reproducible-runtime + PRIVATE rocprofiler-sdk-roctx::rocprofiler-sdk-roctx) + if(REPRODUCIBLE_RUNTIME_USE_MPI) find_package(MPI REQUIRED) target_compile_definitions(reproducible-runtime PRIVATE USE_MPI) diff --git a/tests/apps/reproducible-runtime/reproducible-runtime.cpp b/tests/apps/reproducible-runtime/reproducible-runtime.cpp index a4820085ba..7926c8b312 100644 --- a/tests/apps/reproducible-runtime/reproducible-runtime.cpp +++ b/tests/apps/reproducible-runtime/reproducible-runtime.cpp @@ -21,6 +21,7 @@ // THE SOFTWARE. #include "hip/hip_runtime.h" +#include "rocprofiler-sdk-roctx/roctx.h" #include #include @@ -124,6 +125,8 @@ reproducible_runtime(uint32_t nspin_v) void run(int tid, int devid) { + auto roctx_range_id = roctxRangeStart("run"); + constexpr int min_sa = 8; constexpr int min_avail_simd = 24; dim3 grid(min_sa * min_avail_simd); @@ -140,6 +143,7 @@ run(int tid, int devid) do { + roctxMark("iteration"); uint32_t cyclesleft = 2000 * 1000 * (nruntime - static_cast(time)); HIP_API_CALL(hipEventRecord(start, stream)); reproducible_runtime<<>>(std::min(nspin, cyclesleft)); @@ -166,6 +170,8 @@ run(int tid, int devid) HIP_API_CALL(hipStreamSynchronize(stream)); HIP_API_CALL(hipStreamDestroy(stream)); + + roctxRangeStop(roctx_range_id); } namespace diff --git a/tests/apps/transpose/CMakeLists.txt b/tests/apps/transpose/CMakeLists.txt index 1391f681a2..7835468902 100644 --- a/tests/apps/transpose/CMakeLists.txt +++ b/tests/apps/transpose/CMakeLists.txt @@ -42,6 +42,9 @@ target_compile_options(transpose PRIVATE -W -Wall -Wextra -Wpedantic -Wshadow -W find_package(Threads REQUIRED) target_link_libraries(transpose PRIVATE Threads::Threads) +find_package(rocprofiler-sdk-roctx REQUIRED) +target_link_libraries(transpose PRIVATE rocprofiler-sdk-roctx::rocprofiler-sdk-roctx) + if(TRANSPOSE_USE_MPI) find_package(MPI REQUIRED) target_compile_definitions(transpose PRIVATE USE_MPI) diff --git a/tests/apps/transpose/transpose.cpp b/tests/apps/transpose/transpose.cpp index 2c1b338b17..fb52c8542c 100644 --- a/tests/apps/transpose/transpose.cpp +++ b/tests/apps/transpose/transpose.cpp @@ -21,6 +21,7 @@ // THE SOFTWARE. #include "hip/hip_runtime.h" +#include "rocprofiler-sdk-roctx/roctx.h" #include #include @@ -173,6 +174,8 @@ transpose(const int* in, int* out, int M, int N) void run(int rank, int tid, int devid, int argc, char** argv) { + roctxRangePush("run"); + constexpr unsigned int M = 4960 * 2; constexpr unsigned int N = 4960 * 2; @@ -243,6 +246,8 @@ run(int rank, int tid, int devid, int argc, char** argv) delete[] inp_matrix; delete[] out_matrix; + + roctxRangePop(); } namespace diff --git a/tests/async-copy-tracing/CMakeLists.txt b/tests/async-copy-tracing/CMakeLists.txt index 3f6bdc18ca..a5712ceb02 100644 --- a/tests/async-copy-tracing/CMakeLists.txt +++ b/tests/async-copy-tracing/CMakeLists.txt @@ -19,17 +19,17 @@ endif() add_test(NAME test-async-copy-tracing-execute COMMAND $) +set(async-copy-tracing-env + "${PRELOAD_ENV}" + "HSA_TOOLS_LIB=$" + "ROCPROFILER_TOOL_OUTPUT_FILE=async-copy-tracing-test.json" + "LD_LIBRARY_PATH=$:$ENV{LD_LIBRARY_PATH}" + ) + 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") + PROPERTIES TIMEOUT 45 LABELS "integration-tests" ENVIRONMENT + "${async-copy-tracing-env}" FAIL_REGULAR_EXPRESSION "threw an exception") foreach(FILENAME validate.py pytest.ini conftest.py) configure_file(${CMAKE_CURRENT_SOURCE_DIR}/${FILENAME} diff --git a/tests/async-copy-tracing/validate.py b/tests/async-copy-tracing/validate.py index e0cfa316f6..0838b73eb2 100644 --- a/tests/async-copy-tracing/validate.py +++ b/tests/async-copy-tracing/validate.py @@ -16,45 +16,50 @@ def test_data_structure(input_data): 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"]) + sdk_data = data["rocprofiler-sdk-json-tool"] - 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"]) + node_exists("agents", sdk_data) + node_exists("call_stack", sdk_data) + node_exists("callback_records", sdk_data) + node_exists("buffer_records", sdk_data) + + node_exists("names", sdk_data["callback_records"]) + node_exists("code_objects", sdk_data["callback_records"]) + node_exists("kernel_symbols", sdk_data["callback_records"]) + node_exists("hsa_api_traces", sdk_data["callback_records"]) + + node_exists("names", sdk_data["buffer_records"]) + node_exists("kernel_dispatches", sdk_data["buffer_records"]) + node_exists("memory_copies", sdk_data["buffer_records"], 4) + node_exists("hsa_api_traces", sdk_data["buffer_records"]) def test_timestamps(input_data): data = input_data + sdk_data = data["rocprofiler-sdk-json-tool"] 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 titr in ["hsa_api_traces", "marker_api_traces"]: + for itr in sdk_data["callback_records"][titr]: + 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"]: + for itr in sdk_data["buffer_records"][titr]: + assert itr["start_timestamp"] <= itr["end_timestamp"] + + for itr in sdk_data["buffer_records"]["memory_copies"]: 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"]: + for itr in sdk_data["buffer_records"]["kernel_dispatches"]: assert itr["start_timestamp"] < itr["end_timestamp"] assert itr["correlation_id"]["internal"] > 0 assert itr["correlation_id"]["external"] > 0 @@ -67,21 +72,23 @@ def test_timestamps(input_data): def test_internal_correlation_ids(input_data): data = input_data + sdk_data = data["rocprofiler-sdk-json-tool"] 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 titr in ["hsa_api_traces", "marker_api_traces"]: + for itr in sdk_data["callback_records"][titr]: + 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"]) + for itr in sdk_data["buffer_records"][titr]: + 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"]: + for itr in sdk_data["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"]: + for itr in sdk_data["buffer_records"]["memory_copies"]: assert itr["correlation_id"]["internal"] in api_corr_ids_unique len_corr_id_unq = len(api_corr_ids_unique) @@ -91,34 +98,40 @@ def test_internal_correlation_ids(input_data): def test_external_correlation_ids(input_data): data = input_data + sdk_data = data["rocprofiler-sdk-json-tool"] 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 titr in ["hsa_api_traces", "marker_api_traces"]: + for itr in sdk_data["callback_records"][titr]: + 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 titr in ["hsa_api_traces", "marker_api_traces"]: + for itr in sdk_data["buffer_records"][titr]: + 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"]: + for itr in sdk_data["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"]: + for itr in sdk_data["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 + sdk_data = data["rocprofiler-sdk-json-tool"] symbol_info = {} - for itr in data["rocprofiler-sdk-json-tool"]["callback_records"]["kernel_symbols"]: + for itr in sdk_data["callback_records"]["kernel_symbols"]: phase = itr["record"]["phase"] payload = itr["payload"] kern_id = payload["kernel_id"] @@ -132,12 +145,13 @@ 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["rocprofiler-sdk-json-tool"]["buffer_records"]["kernel_dispatches"]: + for itr in sdk_data["buffer_records"]["kernel_dispatches"]: assert itr["kernel_id"] in symbol_info.keys() def test_async_copy_direction(input_data): data = input_data + sdk_data = data["rocprofiler-sdk-json-tool"] # Direction values: # -1 == ??? (unknown) @@ -146,7 +160,7 @@ def test_async_copy_direction(input_data): # 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"]: + for itr in sdk_data["buffer_records"]["memory_copies"]: op_id = itr["operation"] assert op_id > 0 assert op_id < 3 diff --git a/tests/common/serialization.hpp b/tests/common/serialization.hpp index fa5ce521a1..00484d5abf 100644 --- a/tests/common/serialization.hpp +++ b/tests/common/serialization.hpp @@ -158,6 +158,13 @@ save(ArchiveT& ar, rocprofiler_hsa_api_retval_t data) SAVE_DATA_FIELD(uint64_t_retval); } +template +void +save(ArchiveT& ar, rocprofiler_marker_api_retval_t data) +{ + SAVE_DATA_FIELD(int64_t_retval); +} + template void save(ArchiveT& ar, rocprofiler_callback_tracing_hsa_api_data_t data) @@ -167,6 +174,15 @@ save(ArchiveT& ar, rocprofiler_callback_tracing_hsa_api_data_t data) SAVE_DATA_FIELD(retval); } +template +void +save(ArchiveT& ar, rocprofiler_callback_tracing_marker_api_data_t data) +{ + SAVE_DATA_FIELD(size); + // SAVE_DATA_FIELD(args); + SAVE_DATA_FIELD(retval); +} + template void save(ArchiveT& ar, rocprofiler_callback_tracing_record_t data) diff --git a/tests/kernel-tracing/CMakeLists.txt b/tests/kernel-tracing/CMakeLists.txt index 7483cc25ca..052b9bde07 100644 --- a/tests/kernel-tracing/CMakeLists.txt +++ b/tests/kernel-tracing/CMakeLists.txt @@ -19,17 +19,17 @@ endif() add_test(NAME test-kernel-tracing-execute COMMAND $) +set(kernel-tracing-env + "${PRELOAD_ENV}" + "HSA_TOOLS_LIB=$" + "ROCPROFILER_TOOL_OUTPUT_FILE=kernel-tracing-test.json" + "LD_LIBRARY_PATH=$:$ENV{LD_LIBRARY_PATH}" + ) + set_tests_properties( test-kernel-tracing-execute - PROPERTIES - TIMEOUT - 45 - LABELS - "integration-tests" - ENVIRONMENT - "${PRELOAD_ENV};HSA_TOOLS_LIB=$;ROCPROFILER_TOOL_OUTPUT_FILE=kernel-tracing-test.json" - FAIL_REGULAR_EXPRESSION - "threw an exception") + PROPERTIES TIMEOUT 45 LABELS "integration-tests" ENVIRONMENT "${kernel-tracing-env}" + FAIL_REGULAR_EXPRESSION "threw an exception") foreach(FILENAME validate.py pytest.ini conftest.py) configure_file(${CMAKE_CURRENT_SOURCE_DIR}/${FILENAME} diff --git a/tests/kernel-tracing/validate.py b/tests/kernel-tracing/validate.py index de27cb1a12..58368bbf54 100644 --- a/tests/kernel-tracing/validate.py +++ b/tests/kernel-tracing/validate.py @@ -16,45 +16,52 @@ def test_data_structure(input_data): 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"]) + sdk_data = data["rocprofiler-sdk-json-tool"] - 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"]) + node_exists("agents", sdk_data) + node_exists("call_stack", sdk_data) + node_exists("callback_records", sdk_data) + node_exists("buffer_records", sdk_data) + + node_exists("names", sdk_data["callback_records"]) + node_exists("code_objects", sdk_data["callback_records"]) + node_exists("kernel_symbols", sdk_data["callback_records"]) + node_exists("hsa_api_traces", sdk_data["callback_records"]) + node_exists("marker_api_traces", sdk_data["callback_records"]) + + node_exists("names", sdk_data["buffer_records"]) + node_exists("kernel_dispatches", sdk_data["buffer_records"]) + node_exists("memory_copies", sdk_data["buffer_records"], 0) + node_exists("hsa_api_traces", sdk_data["buffer_records"]) + node_exists("marker_api_traces", sdk_data["buffer_records"]) def test_timestamps(input_data): data = input_data + sdk_data = data["rocprofiler-sdk-json-tool"] 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 titr in ["hsa_api_traces", "marker_api_traces"]: + for itr in sdk_data["callback_records"][titr]: + 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"]: + for itr in sdk_data["buffer_records"][titr]: + assert itr["start_timestamp"] <= itr["end_timestamp"] + + for itr in sdk_data["buffer_records"]["memory_copies"]: 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"]: + for itr in sdk_data["buffer_records"]["kernel_dispatches"]: assert itr["start_timestamp"] < itr["end_timestamp"] assert itr["correlation_id"]["internal"] > 0 assert itr["correlation_id"]["external"] > 0 @@ -67,21 +74,23 @@ def test_timestamps(input_data): def test_internal_correlation_ids(input_data): data = input_data + sdk_data = data["rocprofiler-sdk-json-tool"] 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 titr in ["hsa_api_traces", "marker_api_traces"]: + for itr in sdk_data["callback_records"][titr]: + 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"]) + for itr in sdk_data["buffer_records"][titr]: + 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"]: + for itr in sdk_data["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"]: + for itr in sdk_data["buffer_records"]["memory_copies"]: assert itr["correlation_id"]["internal"] in api_corr_ids_unique len_corr_id_unq = len(api_corr_ids_unique) @@ -91,34 +100,40 @@ def test_internal_correlation_ids(input_data): def test_external_correlation_ids(input_data): data = input_data + sdk_data = data["rocprofiler-sdk-json-tool"] 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 titr in ["hsa_api_traces", "marker_api_traces"]: + for itr in sdk_data["callback_records"][titr]: + 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 titr in ["hsa_api_traces", "marker_api_traces"]: + for itr in sdk_data["buffer_records"][titr]: + 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"]: + for itr in sdk_data["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"]: + for itr in sdk_data["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 + sdk_data = data["rocprofiler-sdk-json-tool"] symbol_info = {} - for itr in data["rocprofiler-sdk-json-tool"]["callback_records"]["kernel_symbols"]: + for itr in sdk_data["callback_records"]["kernel_symbols"]: phase = itr["record"]["phase"] payload = itr["payload"] kern_id = payload["kernel_id"] @@ -132,12 +147,13 @@ 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["rocprofiler-sdk-json-tool"]["buffer_records"]["kernel_dispatches"]: + for itr in sdk_data["buffer_records"]["kernel_dispatches"]: assert itr["kernel_id"] in symbol_info.keys() def test_async_copy_direction(input_data): data = input_data + sdk_data = data["rocprofiler-sdk-json-tool"] # Direction values: # -1 == ??? (unknown) @@ -146,7 +162,7 @@ def test_async_copy_direction(input_data): # 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"]: + for itr in sdk_data["buffer_records"]["memory_copies"]: op_id = itr["operation"] async_dir_cnt[op_id] += 1 diff --git a/tests/tools/CMakeLists.txt b/tests/tools/CMakeLists.txt index b1890890d4..63b6722295 100644 --- a/tests/tools/CMakeLists.txt +++ b/tests/tools/CMakeLists.txt @@ -18,20 +18,13 @@ target_link_libraries( rocprofiler::tests-common-library) set_target_properties( rocprofiler-sdk-json-tool - PROPERTIES INSTALL_RPATH "\$ORIGIN:\$ORIGIN/.." - INSTALL_RPATH_USE_LINK_PATH ON + PROPERTIES LIBRARY_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/lib/rocprofiler-sdk" SOVERSION ${PROJECT_VERSION_MINOR} VERSION - ${PROJECT_VERSION_MAJOR}.${PROJECT_VERSION_MINOR}.${PROJECT_VERSION_PATCH}) + ${PROJECT_VERSION_MAJOR}.${PROJECT_VERSION_MINOR}.${PROJECT_VERSION_PATCH} + INSTALL_RPATH "\$ORIGIN:\$ORIGIN/..") 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/tests/tools/json-tool.cpp b/tests/tools/json-tool.cpp index 32557b827d..ce93f60a08 100644 --- a/tests/tools/json-tool.cpp +++ b/tests/tools/json-tool.cpp @@ -128,6 +128,9 @@ rocprofiler_client_finalize_t client_fini_func = nullptr; callback_name_info get_callback_tracing_names() { + static const auto supported = std::unordered_set{ + ROCPROFILER_CALLBACK_TRACING_HSA_API, ROCPROFILER_CALLBACK_TRACING_MARKER_API}; + auto cb_name_info = callback_name_info{}; // // callback for each kind operation @@ -136,7 +139,7 @@ get_callback_tracing_names() [](rocprofiler_callback_tracing_kind_t kindv, uint32_t operation, void* data_v) { auto* name_info_v = static_cast(data_v); - if(kindv == ROCPROFILER_CALLBACK_TRACING_HSA_API) + if(supported.count(kindv) > 0) { const char* name = nullptr; ROCPROFILER_CALL(rocprofiler_query_callback_tracing_kind_operation_name( @@ -158,7 +161,7 @@ get_callback_tracing_names() "query buffer tracing kind operation name"); if(name) name_info_v->kind_names[kind] = name; - if(kind == ROCPROFILER_CALLBACK_TRACING_HSA_API) + if(supported.count(kind) > 0) { ROCPROFILER_CALL(rocprofiler_iterate_callback_tracing_kind_operations( kind, tracing_kind_operation_cb, static_cast(data)), @@ -177,6 +180,9 @@ get_callback_tracing_names() buffer_name_info get_buffer_tracing_names() { + static const auto supported = std::unordered_set{ + ROCPROFILER_BUFFER_TRACING_HSA_API, ROCPROFILER_BUFFER_TRACING_MARKER_API}; + auto cb_name_info = buffer_name_info{}; // // callback for each kind operation @@ -185,7 +191,7 @@ get_buffer_tracing_names() [](rocprofiler_buffer_tracing_kind_t kindv, uint32_t operation, void* data_v) { auto* name_info_v = static_cast(data_v); - if(kindv == ROCPROFILER_BUFFER_TRACING_HSA_API) + if(supported.count(kindv) > 0) { const char* name = nullptr; ROCPROFILER_CALL(rocprofiler_query_buffer_tracing_kind_operation_name( @@ -207,7 +213,7 @@ get_buffer_tracing_names() "query buffer tracing kind operation name"); if(name) name_info_v->kind_names[kind] = name; - if(kind == ROCPROFILER_BUFFER_TRACING_HSA_API) + if(supported.count(kind) > 0) { ROCPROFILER_CALL(rocprofiler_iterate_buffer_tracing_kind_operations( kind, tracing_kind_operation_cb, static_cast(data)), @@ -226,7 +232,8 @@ get_buffer_tracing_names() using callback_payload_t = std::variant; + rocprofiler_callback_tracing_hsa_api_data_t, + rocprofiler_callback_tracing_marker_api_data_t>; struct code_object_callback_record_t { @@ -273,9 +280,25 @@ struct hsa_api_callback_record_t } }; +struct marker_api_callback_record_t +{ + uint64_t timestamp = 0; + rocprofiler_callback_tracing_record_t record = {}; + rocprofiler_callback_tracing_marker_api_data_t payload = {}; + + template + void save(ArchiveT& ar) const + { + ar(cereal::make_nvp("timestamp", timestamp)); + ar(cereal::make_nvp("record", record)); + ar(cereal::make_nvp("payload", payload)); + } +}; + auto code_object_records = std::deque{}; auto kernel_symbol_records = std::deque{}; auto hsa_api_cb_records = std::deque{}; +auto marker_api_cb_records = std::deque{}; rocprofiler_thread_id_t push_external_correlation(); @@ -303,14 +326,12 @@ tool_tracing_callback(rocprofiler_callback_tracing_record_t record, { auto data_v = *static_cast(record.payload); - data_v.uri = ::strdup(data_v.uri); code_object_records.emplace_back(code_object_callback_record_t{ts, record, data_v}); } else if(record.operation == ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT_DEVICE_KERNEL_SYMBOL_REGISTER) { - auto data_v = *static_cast(record.payload); - data_v.kernel_name = ::strdup(data_v.kernel_name); + auto data_v = *static_cast(record.payload); kernel_symbol_records.emplace_back(kernel_symbol_callback_record_t{ts, record, data_v}); } } @@ -319,6 +340,11 @@ tool_tracing_callback(rocprofiler_callback_tracing_record_t record, auto* data = static_cast(record.payload); hsa_api_cb_records.emplace_back(hsa_api_callback_record_t{ts, record, *data}); } + else if(record.kind == ROCPROFILER_CALLBACK_TRACING_MARKER_API) + { + auto* data = static_cast(record.payload); + marker_api_cb_records.emplace_back(marker_api_callback_record_t{ts, record, *data}); + } else { throw std::runtime_error{"unsupported callback kind"}; @@ -326,6 +352,7 @@ tool_tracing_callback(rocprofiler_callback_tracing_record_t record, } auto hsa_api_bf_records = std::deque{}; +auto marker_api_bf_records = std::deque{}; auto kernel_dispatch_records = std::deque{}; auto memory_copy_records = std::deque{}; @@ -377,6 +404,13 @@ tool_tracing_buffered(rocprofiler_context_id_t /*context*/, hsa_api_bf_records.emplace_back(*record); } + else if(header->kind == ROCPROFILER_BUFFER_TRACING_MARKER_API) + { + auto* record = + static_cast(header->payload); + + marker_api_bf_records.emplace_back(*record); + } else if(header->kind == ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH) { auto* record = static_cast( @@ -444,23 +478,29 @@ void flush(); // contexts -rocprofiler_context_id_t api_callback_ctx = {}; -rocprofiler_context_id_t code_object_ctx = {}; -rocprofiler_context_id_t api_buffered_ctx = {}; -rocprofiler_context_id_t kernel_dispatch_ctx = {}; -rocprofiler_context_id_t memory_copy_ctx = {}; +rocprofiler_context_id_t hsa_api_callback_ctx = {}; +rocprofiler_context_id_t marker_api_callback_ctx = {}; +rocprofiler_context_id_t code_object_ctx = {}; +rocprofiler_context_id_t hsa_api_buffered_ctx = {}; +rocprofiler_context_id_t marker_api_buffered_ctx = {}; +rocprofiler_context_id_t kernel_dispatch_ctx = {}; +rocprofiler_context_id_t memory_copy_ctx = {}; // buffers -rocprofiler_buffer_id_t api_buffered_buffer = {}; -rocprofiler_buffer_id_t kernel_dispatch_buffer = {}; -rocprofiler_buffer_id_t memory_copy_buffer = {}; +rocprofiler_buffer_id_t hsa_api_buffered_buffer = {}; +rocprofiler_buffer_id_t marker_api_buffered_buffer = {}; +rocprofiler_buffer_id_t kernel_dispatch_buffer = {}; +rocprofiler_buffer_id_t memory_copy_buffer = {}; auto contexts = std::unordered_map{ - {"API_CALLBACK", &api_callback_ctx}, + {"HSA_API_CALLBACK", &hsa_api_callback_ctx}, + {"MARKER_API_CALLBACK", &marker_api_callback_ctx}, {"CODE_OBJECT", &code_object_ctx}, - {"API_BUFFERED", &api_buffered_ctx}, + {"HSA_API_BUFFERED", &hsa_api_buffered_ctx}, + {"MARKER_API_BUFFERED", &marker_api_buffered_ctx}, {"KERNEL_DISPATCH", &kernel_dispatch_ctx}, {"MEMORY_COPY", &memory_copy_ctx}}; -auto buffers = std::array{&api_buffered_buffer, +auto buffers = std::array{&hsa_api_buffered_buffer, + &marker_api_buffered_buffer, &kernel_dispatch_buffer, &memory_copy_buffer}; @@ -475,14 +515,7 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data) [](const rocprofiler_agent_t** agents_arr, size_t num_agents, void* user_data) { auto* agents_v = static_cast*>(user_data); for(size_t i = 0; i < num_agents; ++i) - { - const auto* agent = agents_arr[i]; - auto& val = agents_v->emplace_back(*agent); - val.name = ::strdup(agent->name); - val.vendor_name = ::strdup(agent->vendor_name); - val.product_name = ::strdup(agent->product_name); - val.model_name = ::strdup(agent->model_name); - } + agents_v->emplace_back(*agents_arr[i]); return ROCPROFILER_STATUS_SUCCESS; }; @@ -504,7 +537,7 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data) } ROCPROFILER_CALL( - rocprofiler_configure_callback_tracing_service(api_callback_ctx, + rocprofiler_configure_callback_tracing_service(hsa_api_callback_ctx, ROCPROFILER_CALLBACK_TRACING_HSA_API, nullptr, 0, @@ -521,16 +554,34 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data) nullptr), "code object tracing service configure"); + ROCPROFILER_CALL( + rocprofiler_configure_callback_tracing_service(marker_api_callback_ctx, + ROCPROFILER_CALLBACK_TRACING_MARKER_API, + nullptr, + 0, + tool_tracing_callback, + nullptr), + "hsa api tracing service configure"); + constexpr auto buffer_size = 8192; constexpr auto watermark = 7936; - ROCPROFILER_CALL(rocprofiler_create_buffer(api_buffered_ctx, + ROCPROFILER_CALL(rocprofiler_create_buffer(hsa_api_buffered_ctx, buffer_size, watermark, ROCPROFILER_BUFFER_POLICY_LOSSLESS, tool_tracing_buffered, tool_data, - &api_buffered_buffer), + &hsa_api_buffered_buffer), + "buffer creation"); + + ROCPROFILER_CALL(rocprofiler_create_buffer(marker_api_buffered_ctx, + buffer_size, + watermark, + ROCPROFILER_BUFFER_POLICY_LOSSLESS, + tool_tracing_buffered, + tool_data, + &marker_api_buffered_buffer), "buffer creation"); ROCPROFILER_CALL(rocprofiler_create_buffer(kernel_dispatch_ctx, @@ -552,8 +603,19 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data) "buffer creation"); ROCPROFILER_CALL( - rocprofiler_configure_buffer_tracing_service( - api_buffered_ctx, ROCPROFILER_BUFFER_TRACING_HSA_API, nullptr, 0, api_buffered_buffer), + rocprofiler_configure_buffer_tracing_service(hsa_api_buffered_ctx, + ROCPROFILER_BUFFER_TRACING_HSA_API, + nullptr, + 0, + hsa_api_buffered_buffer), + "buffer tracing service configure"); + + ROCPROFILER_CALL( + rocprofiler_configure_buffer_tracing_service(marker_api_buffered_ctx, + ROCPROFILER_BUFFER_TRACING_MARKER_API, + nullptr, + 0, + marker_api_buffered_buffer), "buffer tracing service configure"); ROCPROFILER_CALL( @@ -650,9 +712,11 @@ tool_fini(void* tool_data) << ", code_object_callback_records=" << code_object_records.size() << ", kernel_symbol_callback_records=" << kernel_symbol_records.size() << ", hsa_api_callback_records=" << hsa_api_cb_records.size() + << ", marker_api_callback_records=" << marker_api_cb_records.size() << ", kernel_dispatch_records=" << kernel_dispatch_records.size() << ", memory_copy_records=" << memory_copy_records.size() - << ", hsa_api_bf_records=" << hsa_api_bf_records.size() << " ...\n" + << ", hsa_api_bf_records=" << hsa_api_bf_records.size() + << ", marker_api_bf_records=" << marker_api_bf_records.size() << " ...\n" << std::flush; auto* _call_stack = static_cast(tool_data); @@ -713,6 +777,7 @@ tool_fini(void* tool_data) json_ar(cereal::make_nvp("code_objects", code_object_records)); json_ar(cereal::make_nvp("kernel_symbols", kernel_symbol_records)); json_ar(cereal::make_nvp("hsa_api_traces", hsa_api_cb_records)); + json_ar(cereal::make_nvp("marker_api_traces", marker_api_cb_records)); } catch(std::exception& e) { std::cerr << "[" << getpid() << "][" << __FUNCTION__ @@ -729,6 +794,7 @@ tool_fini(void* tool_data) json_ar(cereal::make_nvp("kernel_dispatches", kernel_dispatch_records)); json_ar(cereal::make_nvp("memory_copies", memory_copy_records)); json_ar(cereal::make_nvp("hsa_api_traces", hsa_api_bf_records)); + json_ar(cereal::make_nvp("marker_api_traces", marker_api_bf_records)); } catch(std::exception& e) { std::cerr << "[" << getpid() << "][" << __FUNCTION__ @@ -748,26 +814,6 @@ tool_fini(void* tool_data) << std::flush; delete _call_stack; - - auto free_cstr = [](const char*& val) { - ::free(const_cast(val)); - val = nullptr; - }; - - // clean up our strdups to avoid triggering our leak sanitizer during CI - for(auto& itr : code_object_records) - free_cstr(itr.payload.uri); - - for(auto& itr : kernel_symbol_records) - free_cstr(itr.payload.kernel_name); - - for(auto& itr : agents) - { - free_cstr(itr.name); - free_cstr(itr.vendor_name); - free_cstr(itr.product_name); - free_cstr(itr.model_name); - } } void