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

ROCTx Library Tracing (#390)

Browse Source 
* Update include/rocprofiler-sdk/marker/*

- Update rocprofiler_marker_api_args_t for all API functions
- Add ROCPROFILER_MARKER_API_ID_roctxGetThreadId to rocprofiler_marker_api_id_t

* Update include/rocprofiler-sdk/marker/api_args.h

- fix include

* Update lib/common/mpl.hpp

- is_pair
- is_type_complete_v

* Update include/rocprofiler-sdk/marker/*

- fix rocprofiler_marker_api_retval_t
- add roctxGetThreadId to rocprofiler_marker_api_args_t
- fix type in enum: HsaDevice -> HsaAgent
- add table_api_id.h

* Update include/rocprofiler-sdk/marker.h

- include marker/table_api_id.h

* Update include/rocprofiler-sdk/buffer_tracing.h

- Buffer marker tracer records have begin and end timestamp

* Add lib/rocprofiler-sdk/marker

- tracing implementation for marker (roctx) library

* Update include/rocprofiler-sdk/{buffer_tracing,marker/table_api_id}.h

- rocprofiler_buffer_tracing_marker_record_t -> rocprofiler_buffer_tracing_marker_api_record_t

* Update lib/rocprofiler-sdk/buffer_tracing.cpp

- support for ROCPROFILER_BUFFER_TRACING_MARKER_API

* Update lib/rocprofiler-sdk/callback_tracing.cpp

- support for ROCPROFILER_CALLBACK_TRACING_MARKER_API

* Update lib/rocprofiler-sdk/intercept_table.cpp

- template instantiation for notify_runtime_api_registration

* Update lib/rocprofiler-sdk/registration.cpp

- enable roctx in rocprofiler_set_api_table

* Update lib/rocprofiler-sdk/marker/marker.cpp

- rocprofiler_buffer_tracing_marker_record_t -> rocprofiler_buffer_tracing_marker_api_record_t

* Update lib/rocprofiler/tests for roctx testing

- add roctx.cpp
  - unit tests for roctx callback and buffer tracing
- support marker API in get_{buffer,callback}_tracing_names()

* Update lib/common/logging.cpp

- logging initialized message mentions env variable

* Update lib/common/mpl.hpp

- NOLINT for misc-definitions-in-headers

* Update lib/rocprofiler-sdk/tests/CMakeLists.txt

- include LD_LIBRARY_PATH in rocprofiler-lib-tests-shared tests

* Update lib/rocprofiler-sdk/registration.cpp

- client_library_vec_t is now vector of option<client_library>
  - enables resetting the client_library after finalization
- removed acquiring registration lock when invoke_client_finalizers called via atexit
  - this was causing some lock-order-inversion warnings (potential deadlock)

* Update lib/rocprofiler-sdk/agent.cpp

- model name for agent supports spaces

* Update tests/common/serialization.hpp

- add serialization support for marker tracing data structures

* Update tests/apps

- Add ROCTx markers into reproducible-runtime and transpose

* Update tests/tools/json-tools.cpp

- add marker tracing support
- remove strdup (no longer necessary)

* Update tests/kernel-tracing/validate.py

- validate marker API tracing data

* Update tests/async-copy-tracing/validate.py

- validate marker API tracing data

* Update cmake for load path resolution during testing

* Update tests/async-copy-tracing/CMakeLists.txt

- fix test LD_LIBRARY_PATH

* Update cmake/Templates/rocprofiler-sdk-roctx/config.cmake.in

- fix constructing rocprofiler-sdk-roctx::rocprofiler-sdk-roctx
This commit is contained in:
Jonathan R. Madsen
2024-01-18 09:48:06 -06:00
committed by GitHub
parent b55cea0e98
commit 21dd088c8e
37 changed files with 2270 additions and 310 deletions
Download Patch File Download Diff File Expand all files Collapse all files
cmake/Templates/rocprofiler-sdk-roctx/config.cmake.in
+2 -2
View File
@@ -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})
source/include/rocprofiler-sdk/buffer_tracing.h
+6 -5
View File
@@ -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;
source/include/rocprofiler-sdk/marker.h
+1
View File
@@ -24,3 +24,4 @@
#include <rocprofiler-sdk/marker/api_args.h>
#include <rocprofiler-sdk/marker/api_id.h>
#include <rocprofiler-sdk/marker/table_api_id.h>
source/include/rocprofiler-sdk/marker/CMakeLists.txt
+1 -1
View File
@@ -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}
source/include/rocprofiler-sdk/marker/api_args.h
+36 -3
View File
@@ -22,13 +22,15 @@
#pragma once
#include <stdint.h>
#include <rocprofiler-sdk-roctx/api_trace.h>
#include <rocprofiler-sdk-roctx/types.h>
typedef uint64_t roctx_range_id_t;
#include <stdint.h>
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;
source/include/rocprofiler-sdk/marker/api_id.h
+2 -1
View File
@@ -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;
source/include/rocprofiler-sdk/marker/table_api_id.h
+30
View File
@@ -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;
source/lib/common/logging.cpp
+1 -1
View File
@@ -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
source/lib/common/mpl.hpp
+22
View File
@@ -90,6 +90,28 @@ struct is_one_of<F, Tuple<S, T...>>
{
static constexpr bool value = is_one_of<F, S, Tuple<T...>>::value;
};
template <typename Tp>
struct is_pair_impl
{
static constexpr auto value = false;
};
template <typename LhsT, typename RhsT>
struct is_pair_impl<std::pair<LhsT, RhsT>>
{
static constexpr auto value = true;
};
template <typename Tp>
struct is_pair : is_pair_impl<std::remove_cv_t<std::remove_reference_t<std::decay_t<Tp>>>>
{};
template <typename, typename = void>
constexpr bool is_type_complete_v = false; // NOLINT(misc-definitions-in-headers)
template <typename T> // NOLINTNEXTLINE(misc-definitions-in-headers)
constexpr bool is_type_complete_v<T, std::void_t<decltype(sizeof(T))>> = true;
} // namespace mpl
} // namespace common
} // namespace rocprofiler
source/lib/rocprofiler-sdk-roctx/CMakeLists.txt
+2
View File
@@ -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(
source/lib/rocprofiler-sdk/CMakeLists.txt
+1
View File
@@ -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
source/lib/rocprofiler-sdk/agent.cpp
+3 -2
View File
@@ -30,7 +30,7 @@
#include "lib/rocprofiler-sdk/agent.hpp"
#include "lib/rocprofiler-sdk/hsa/agent_cache.hpp"
#include <fmt/core.h>
#include <fmt/format.h>
#include <glog/logging.h>
#include <hsa/hsa_api_trace.h>
#include <libdrm/amdgpu.h>
@@ -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 = "";
source/lib/rocprofiler-sdk/buffer_tracing.cpp
+28 -24
View File
@@ -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 <glog/logging.h>
@@ -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<ROCPROFILER_MARKER_API_TABLE_ID_RoctxApi>(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<uint32_t>{};
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<ROCPROFILER_MARKER_API_TABLE_ID_RoctxApi>();
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;
}
}
source/lib/rocprofiler-sdk/callback_tracing.cpp
+37 -24
View File
@@ -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 <glog/logging.h>
@@ -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<ROCPROFILER_MARKER_API_TABLE_ID_RoctxApi>(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<uint32_t>{};
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<ROCPROFILER_MARKER_API_TABLE_ID_RoctxApi>();
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<rocprofiler_callback_tracing_marker_api_data_t*>(record.payload),
callback,
user_data);
return ROCPROFILER_STATUS_SUCCESS;
}
return ROCPROFILER_STATUS_ERROR_NOT_IMPLEMENTED;
}
source/lib/rocprofiler-sdk/intercept_table.cpp
+5
View File
@@ -156,6 +156,11 @@ template void notify_runtime_api_registration(rocprofiler_runtime_library_t,
uint64_t,
uint64_t,
std::tuple<HsaApiTable*>);
template void notify_runtime_api_registration(rocprofiler_runtime_library_t,
uint64_t,
uint64_t,
std::tuple<roctxApiTable_t*>);
} // namespace intercept_table
} // namespace rocprofiler
source/lib/rocprofiler-sdk/marker/CMakeLists.txt
+7
View File
@@ -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})
source/lib/rocprofiler-sdk/marker/defines.hpp
+286
View File
@@ -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<void*>(&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<MARKER_TABLE, MARKER_API_ID> : roctx_domain_info<MARKER_TABLE> \
{ \
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<table_idx>; \
using this_type = roctx_api_info<table_idx, operation_idx>; \
using base_type = roctx_api_impl<table_idx, operation_idx>; \
\
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<table_idx>::type, MARKER_FUNC_PTR); \
} \
\
static_assert(offsetof(roctx_table_lookup<table_idx>::type, MARKER_FUNC_PTR) == \
(sizeof(size_t) + (operation_idx * sizeof(void*))), \
"ABI error for " #MARKER_FUNC); \
\
static auto& get_table() { return roctx_table_lookup<table_idx>{}(); } \
\
template <typename TableT> \
static auto& get_table(TableT& _v) \
{ \
return roctx_table_lookup<table_idx>{}(_v); \
} \
\
template <typename TableT> \
static auto& get_table_func(TableT& _table) \
{ \
if constexpr(std::is_pointer<TableT>::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 <typename DataT> \
static auto& get_api_data_args(DataT& _data) \
{ \
return _data.MARKER_FUNC; \
} \
\
template <typename RetT, typename... Args> \
static auto get_functor(RetT (*)(Args...)) \
{ \
if constexpr(std::is_void<RetT>::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<void*> as_arg_addr(callback_data_type) { return std::vector<void*>{}; } \
\
static std::vector<std::pair<std::string, std::string>> 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<MARKER_TABLE, MARKER_API_ID> : roctx_domain_info<MARKER_TABLE> \
{ \
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<table_idx>; \
using this_type = roctx_api_info<table_idx, operation_idx>; \
using base_type = roctx_api_impl<table_idx, operation_idx>; \
\
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<table_idx>::type, MARKER_FUNC_PTR); \
} \
\
static_assert(offsetof(roctx_table_lookup<table_idx>::type, MARKER_FUNC_PTR) == \
(sizeof(size_t) + (operation_idx * sizeof(void*))), \
"ABI error for " #MARKER_FUNC); \
\
static auto& get_table() { return roctx_table_lookup<table_idx>{}(); } \
\
template <typename TableT> \
static auto& get_table(TableT& _v) \
{ \
return roctx_table_lookup<table_idx>{}(_v); \
} \
\
template <typename TableT> \
static auto& get_table_func(TableT& _table) \
{ \
if constexpr(std::is_pointer<TableT>::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 <typename DataT> \
static auto& get_api_data_args(DataT& _data) \
{ \
return _data.MARKER_FUNC; \
} \
\
template <typename RetT, typename... Args> \
static auto get_functor(RetT (*)(Args...)) \
{ \
if constexpr(std::is_same<RetT, 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<void*> as_arg_addr(callback_data_type trace_data) \
{ \
return std::vector<void*>{ \
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<TABLE_ID> \
{ \
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()); } \
}; \
} \
}
source/lib/rocprofiler-sdk/marker/marker.cpp
+612
View File
@@ -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 <rocprofiler-sdk/buffer.h>
#include <rocprofiler-sdk/callback_tracing.h>
#include <rocprofiler-sdk/fwd.h>
#include <rocprofiler-sdk/marker.h>
#include <glog/logging.h>
#include <rocprofiler-sdk-roctx/roctx.h>
#include <atomic>
#include <cstddef>
#include <cstdint>
#include <type_traits>
#include <utility>
namespace rocprofiler
{
namespace marker
{
namespace
{
struct null_type
{};
template <typename Tp>
auto
get_default_retval()
{
if constexpr(std::is_integral<Tp>::value)
return Tp{0};
else
static_assert(std::is_empty<Tp>::value, "Error! unsupported return type");
}
template <typename DataT, typename Tp>
void
set_data_retval(DataT& _data, Tp _val)
{
if constexpr(std::is_same<int32_t, Tp>::value)
_data.int32_t_retval = _val;
else if constexpr(std::is_same<int64_t, Tp>::value)
_data.int64_t_retval = _val;
else if constexpr(std::is_same<uint64_t, Tp>::value)
_data.uint64_t_retval = _val;
else
static_assert(std::is_empty<Tp>::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 <size_t TableIdx, size_t OpIdx>
template <typename DataArgsT, typename... Args>
auto
roctx_api_impl<TableIdx, OpIdx>::set_data_args(DataArgsT& _data_args, Args... args)
{
_data_args = DataArgsT{args...};
}
template <size_t TableIdx, size_t OpIdx>
template <typename FuncT, typename... Args>
auto
roctx_api_impl<TableIdx, OpIdx>::exec(FuncT&& _func, Args&&... args)
{
using return_type = std::decay_t<std::invoke_result_t<FuncT, Args...>>;
if(_func)
{
if constexpr(std::is_void<return_type>::value)
{
_func(std::forward<Args>(args)...);
return null_type{};
}
else
{
return _func(std::forward<Args>(args)...);
}
}
using info_type = roctx_api_info<TableIdx, OpIdx>;
LOG(ERROR) << "nullptr to next roctx function for " << info_type::name << " ("
<< info_type::operation_idx << ")";
if constexpr(std::is_void<return_type>::value)
return null_type{};
else
return get_default_retval<return_type>();
}
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_context_data>& callback_contexts,
std::vector<buffered_context_data>& 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 <size_t TableIdx, size_t OpIdx>
template <typename... Args>
auto
roctx_api_impl<TableIdx, OpIdx>::functor(Args&&... args)
{
using info_type = roctx_api_info<TableIdx, OpIdx>;
using callback_api_data_t = typename roctx_domain_info<TableIdx>::callback_data_type;
auto thr_id = common::get_tid();
auto callback_contexts = std::vector<callback_context_data>{};
auto buffered_contexts = std::vector<buffered_context_data>{};
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>(args)...);
if constexpr(!std::is_same<decltype(_ret), null_type>::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>(args)...);
// invoke the callbacks
if(!callback_contexts.empty())
{
set_data_args(info_type::get_api_data_args(tracer_data.args), std::forward<Args>(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<void*>(&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>(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<void*>(&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<decltype(_ret), null_type>::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 <size_t TableIdx>
struct api_id_bounds;
template <>
struct api_id_bounds<ROCPROFILER_MARKER_API_TABLE_ID_RoctxApi>
{
static constexpr auto none = ROCPROFILER_MARKER_API_ID_NONE;
static constexpr auto last = ROCPROFILER_MARKER_API_ID_LAST;
};
template <size_t TableIdx, size_t OpIdx, size_t... OpIdxTail>
const char*
name_by_id(const uint32_t id, std::index_sequence<OpIdx, OpIdxTail...>)
{
if(OpIdx == id) return roctx_api_info<TableIdx, OpIdx>::name;
if constexpr(sizeof...(OpIdxTail) > 0)
return name_by_id<TableIdx>(id, std::index_sequence<OpIdxTail...>{});
else
return nullptr;
}
template <size_t TableIdx, size_t OpIdx, size_t... OpIdxTail>
uint32_t
id_by_name(const char* name, std::index_sequence<OpIdx, OpIdxTail...>)
{
if(std::string_view{roctx_api_info<TableIdx, OpIdx>::name} == std::string_view{name})
return roctx_api_info<TableIdx, OpIdx>::operation_idx;
if constexpr(sizeof...(OpIdxTail) > 0)
return id_by_name<TableIdx>(name, std::index_sequence<OpIdxTail...>{});
else
return api_id_bounds<TableIdx>::none;
}
template <size_t TableIdx, size_t OpIdx, size_t... OpIdxTail>
void
get_ids(std::vector<uint32_t>& _id_list, std::index_sequence<OpIdx, OpIdxTail...>)
{
auto _idx = roctx_api_info<TableIdx, OpIdx>::operation_idx;
if(_idx < api_id_bounds<TableIdx>::last) _id_list.emplace_back(_idx);
if constexpr(sizeof...(OpIdxTail) > 0)
get_ids<TableIdx>(_id_list, std::index_sequence<OpIdxTail...>{});
}
template <size_t TableIdx, size_t OpIdx, size_t... OpIdxTail>
void
get_names(std::vector<const char*>& _name_list, std::index_sequence<OpIdx, OpIdxTail...>)
{
auto&& _name = roctx_api_info<TableIdx, OpIdx>::name;
if(_name != nullptr && strnlen(_name, 1) > 0) _name_list.emplace_back(_name);
if constexpr(sizeof...(OpIdxTail) > 0)
get_names<TableIdx>(_name_list, std::index_sequence<OpIdxTail...>{});
}
template <size_t OpIdx, size_t... OpIdxTail>
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<OpIdx, OpIdxTail...>)
{
if(OpIdx == id)
{
using info_type = roctx_api_info<ROCPROFILER_MARKER_API_TABLE_ID_RoctxApi, OpIdx>;
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<OpIdxTail...>{});
}
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 <size_t TableIdx, typename Tp, size_t OpIdx>
void
copy_table(Tp* _orig, std::integral_constant<size_t, OpIdx>)
{
using table_type = typename roctx_table_lookup<TableIdx>::type;
if constexpr(std::is_same<table_type, Tp>::value)
{
auto _info = roctx_api_info<TableIdx, OpIdx>{};
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 <size_t TableIdx, typename Tp, size_t OpIdx>
void
update_table(Tp* _orig, std::integral_constant<size_t, OpIdx>)
{
using table_type = typename roctx_table_lookup<TableIdx>::type;
if constexpr(std::is_same<table_type, Tp>::value)
{
auto _info = roctx_api_info<TableIdx, OpIdx>{};
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 <size_t TableIdx, typename Tp, size_t OpIdx, size_t... OpIdxTail>
void
copy_table(Tp* _orig, std::index_sequence<OpIdx, OpIdxTail...>)
{
copy_table<TableIdx>(_orig, std::integral_constant<size_t, OpIdx>{});
if constexpr(sizeof...(OpIdxTail) > 0)
copy_table<TableIdx>(_orig, std::index_sequence<OpIdxTail...>{});
}
template <size_t TableIdx, typename Tp, size_t OpIdx, size_t... OpIdxTail>
void
update_table(Tp* _orig, std::index_sequence<OpIdx, OpIdxTail...>)
{
update_table<TableIdx>(_orig, std::integral_constant<size_t, OpIdx>{});
if constexpr(sizeof...(OpIdxTail) > 0)
update_table<TableIdx>(_orig, std::index_sequence<OpIdxTail...>{});
}
} // namespace
// check out the assembly here... this compiles to a switch statement
template <size_t TableIdx>
const char*
name_by_id(uint32_t id)
{
return name_by_id<TableIdx>(id, std::make_index_sequence<api_id_bounds<TableIdx>::last>{});
}
template <size_t TableIdx>
uint32_t
id_by_name(const char* name)
{
return id_by_name<TableIdx>(name, std::make_index_sequence<api_id_bounds<TableIdx>::last>{});
}
template <size_t TableIdx>
std::vector<uint32_t>
get_ids()
{
constexpr auto last_api_id = api_id_bounds<TableIdx>::last;
auto _data = std::vector<uint32_t>{};
_data.reserve(last_api_id);
get_ids<TableIdx>(_data, std::make_index_sequence<last_api_id>{});
return _data;
}
template <size_t TableIdx>
std::vector<const char*>
get_names()
{
constexpr auto last_api_id = api_id_bounds<TableIdx>::last;
auto _data = std::vector<const char*>{};
_data.reserve(last_api_id);
get_names<TableIdx>(_data, std::make_index_sequence<last_api_id>{});
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<ROCPROFILER_MARKER_API_ID_LAST>{});
}
#define INSTANTIATE_MARKER_TABLE_FUNC(TABLE) \
template const char* name_by_id<TABLE>(uint32_t); \
template uint32_t id_by_name<TABLE>(const char*); \
template std::vector<uint32_t> get_ids<TABLE>(); \
template std::vector<const char*> get_names<TABLE>();
INSTANTIATE_MARKER_TABLE_FUNC(ROCPROFILER_MARKER_API_TABLE_ID_RoctxApi)
void
copy_table(roctx_api_table_t* _orig)
{
if(_orig)
copy_table<ROCPROFILER_MARKER_API_TABLE_ID_RoctxApi>(
_orig, std::make_index_sequence<ROCPROFILER_MARKER_API_ID_LAST>{});
}
void
update_table(roctx_api_table_t* _orig)
{
if(_orig)
update_table<ROCPROFILER_MARKER_API_TABLE_ID_RoctxApi>(
_orig, std::make_index_sequence<ROCPROFILER_MARKER_API_ID_LAST>{});
}
} // namespace marker
} // namespace rocprofiler
source/lib/rocprofiler-sdk/marker/marker.def.cpp
+65
View File
@@ -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<ROCPROFILER_MARKER_API_TABLE_ID_RoctxApi>
{
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
source/lib/rocprofiler-sdk/marker/marker.hpp
+91
View File
@@ -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 <rocprofiler-sdk/rocprofiler.h>
#include <rocprofiler-sdk-roctx/api_trace.h>
#include <cstdint>
#include <vector>
namespace rocprofiler
{
namespace marker
{
using roctx_api_table_t = ::roctxApiTable_t;
roctx_api_table_t&
get_table();
template <size_t OpIdx>
struct roctx_table_lookup;
template <size_t TableIdx>
struct roctx_domain_info;
template <size_t TableIdx, size_t OpIdx>
struct roctx_api_info;
template <size_t TableIdx, size_t OpIdx>
struct roctx_api_impl : roctx_domain_info<TableIdx>
{
template <typename DataArgsT, typename... Args>
static auto set_data_args(DataArgsT&, Args... args);
template <typename FuncT, typename... Args>
static auto exec(FuncT&&, Args&&... args);
template <typename... Args>
static auto functor(Args&&... args);
};
template <size_t TableIdx>
const char*
name_by_id(uint32_t id);
template <size_t TableIdx>
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 <size_t TableIdx>
std::vector<const char*>
get_names();
template <size_t TableIdx>
std::vector<uint32_t>
get_ids();
void
copy_table(roctx_api_table_t* _orig);
void
update_table(roctx_api_table_t* _orig);
} // namespace marker
} // namespace rocprofiler
source/lib/rocprofiler-sdk/marker/utils.hpp
+124
View File
@@ -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 <hsa/hsa.h>
#include <rocprofiler-sdk/version.h>
#include "lib/common/mpl.hpp"
#include "fmt/core.h"
#include "fmt/ranges.h"
#include <sstream>
#include <string>
#include <string_view>
#include <type_traits>
#include <utility>
#include <vector>
namespace rocprofiler
{
namespace marker
{
namespace utils
{
template <typename Tp>
auto
stringize_impl(const Tp& _v)
{
using nonpointer_type = typename std::remove_pointer_t<Tp>;
if constexpr(common::mpl::is_pair<Tp>::value)
{
return std::make_pair(stringize_impl(_v.first), stringize_impl(_v.second));
}
else if constexpr(std::is_constructible<std::string_view, Tp>::value)
{
auto _ss = std::stringstream{};
_ss << _v;
return _ss.str();
}
else if constexpr(fmt::is_formattable<Tp>::value && !std::is_pointer<Tp>::value)
{
return fmt::format("{}", _v);
}
else if constexpr(std::is_pointer<Tp>::value && !std::is_pointer<nonpointer_type>::value &&
common::mpl::is_type_complete_v<nonpointer_type> &&
!std::is_void<nonpointer_type>::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 <typename... Args>
auto
stringize(Args... args)
{
return std::vector<std::pair<std::string, std::string>>{stringize_impl(args)...};
}
template <typename Tp>
struct handle_formatter
{
template <typename ParseContext>
constexpr auto parse(ParseContext& ctx)
{
return ctx.begin();
}
template <typename Ctx>
auto format(const Tp& v, Ctx& ctx) const
{
return fmt::format_to(ctx.out(), "handle={}", v.handle);
}
};
template <typename Tp>
struct handle_formatter<const Tp> : handle_formatter<Tp>
{};
} // namespace utils
} // namespace marker
} // namespace rocprofiler
namespace fmt
{
template <>
struct formatter<hsa_agent_t> : rocprofiler::marker::utils::handle_formatter<hsa_agent_t>
{};
} // namespace fmt
source/lib/rocprofiler-sdk/registration.cpp
+66 -63
View File
@@ -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 <rocprofiler-sdk/context.h>
#include <rocprofiler-sdk/fwd.h>
@@ -156,7 +157,7 @@ struct client_library
rocprofiler_client_id_t mutable_client_id = {};
};
using client_library_vec_t = std::vector<client_library>;
using client_library_vec_t = std::vector<std::optional<client_library>>;
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>;
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<uint64_t>(itr.configure_func))
<< itr->name << " (addr="
<< fmt::format("{:#018x}", reinterpret_cast<uint64_t>(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<uint64_t>(itr.configure_func))
<< itr->name << " (addr="
<< fmt::format("{:#018x}", reinterpret_cast<uint64_t>(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<roctxApiTable_t*>(*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
{
source/lib/rocprofiler-sdk/tests/CMakeLists.txt
+14 -6
View File
@@ -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=$<TARGET_FILE:rocprofiler-shared-library>"
"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=$<TARGET_FILE:rocprofiler-shared-library>")
${shared_lib_TESTS} PROPERTIES TIMEOUT 360 LABELS "unittests" ENVIRONMENT
"${rocprofiler-lib-tests-env}")
source/lib/rocprofiler-sdk/tests/common.hpp
+4 -3
View File
@@ -25,6 +25,7 @@
#include <rocprofiler-sdk/rocprofiler.h>
#include "lib/common/defines.hpp"
#include "rocprofiler-sdk/fwd.h"
#include <gtest/gtest.h>
@@ -76,7 +77,7 @@ inline auto
get_callback_tracing_names()
{
static const auto supported_kinds = std::unordered_set<rocprofiler_callback_tracing_kind_t>{
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_kind_t>{ROCPROFILER_BUFFER_TRACING_HSA_API};
static const auto supported_kinds = std::unordered_set<rocprofiler_buffer_tracing_kind_t>{
ROCPROFILER_BUFFER_TRACING_HSA_API, ROCPROFILER_BUFFER_TRACING_MARKER_API};
auto cb_name_info = buffer_name_info{};
//
source/lib/rocprofiler-sdk/tests/roctx.cpp
+542
View File
@@ -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 <hsa/hsa.h>
#include <rocprofiler-sdk-roctx/roctx.h>
#include <rocprofiler-sdk/buffer.h>
#include <rocprofiler-sdk/fwd.h>
#include <rocprofiler-sdk/registration.h>
#include <rocprofiler-sdk/rocprofiler.h>
#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 <gtest/gtest.h>
#include <dlfcn.h>
#include <pthread.h>
#include <chrono>
#include <cmath>
#include <cstdint>
#include <cstdlib>
#include <iostream>
#include <random>
#include <sstream>
#include <string>
#include <string_view>
#include <typeinfo>
#include <unordered_map>
#include <vector>
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<callback_data*>(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<callback_data*>(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<info_data*>(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<void*>(&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<callback_data*>(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<rocprofiler_buffer_tracing_marker_api_record_t*>{};
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<rocprofiler_buffer_tracing_marker_api_record_t*>(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<int64_t>(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<callback_data*>(tool_data);
cb_data->client_workflow_count++;
}
void
thread_postcreate(rocprofiler_runtime_library_t /*lib*/, void* tool_data)
{
auto* cb_data = static_cast<callback_data*>(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<callback_data*>(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<uint32_t>{};
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<std::vector<uint32_t>*>(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<uint32_t, 2>{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<callback_data*>(client_data);
ROCPROFILER_CALL(rocprofiler_stop_context(cb_data->client_ctx),
"rocprofiler context stop failed");
static_cast<callback_data*>(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<void*>(&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<callback_data*>(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<uint32_t, 2>{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<uint32_t>{};
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<std::vector<uint32_t>*>(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<callback_data*>(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<callback_data*>(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<void*>(&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<void*>(&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);
}
tests/apps/CMakeLists.txt
+4
View File
@@ -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:$<TARGET_FILE_DIR:rocprofiler-sdk-roctx::rocprofiler-sdk-roctx-shared-library>"
)
# applications used by integration tests
add_subdirectory(reproducible-runtime)
add_subdirectory(transpose)
tests/apps/reproducible-runtime/CMakeLists.txt
+4
View File
@@ -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)
tests/apps/reproducible-runtime/reproducible-runtime.cpp
+6
View File
@@ -21,6 +21,7 @@
// THE SOFTWARE.
#include "hip/hip_runtime.h"
#include "rocprofiler-sdk-roctx/roctx.h"
#include <unistd.h>
#include <chrono>
@@ -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<double>(time));
HIP_API_CALL(hipEventRecord(start, stream));
reproducible_runtime<<<grid, block, 0, stream>>>(std::min<uint32_t>(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
tests/apps/transpose/CMakeLists.txt
+3
View File
@@ -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)
tests/apps/transpose/transpose.cpp
+5
View File
@@ -21,6 +21,7 @@
// THE SOFTWARE.
#include "hip/hip_runtime.h"
#include "rocprofiler-sdk-roctx/roctx.h"
#include <chrono>
#include <cstdio>
@@ -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
tests/async-copy-tracing/CMakeLists.txt
+9 -9
View File
@@ -19,17 +19,17 @@ endif()
add_test(NAME test-async-copy-tracing-execute COMMAND $<TARGET_FILE:transpose>)
set(async-copy-tracing-env
"${PRELOAD_ENV}"
"HSA_TOOLS_LIB=$<TARGET_FILE:rocprofiler::rocprofiler-shared-library>"
"ROCPROFILER_TOOL_OUTPUT_FILE=async-copy-tracing-test.json"
"LD_LIBRARY_PATH=$<TARGET_FILE_DIR:rocprofiler::rocprofiler-shared-library>:$ENV{LD_LIBRARY_PATH}"
)
set_tests_properties(
test-async-copy-tracing-execute
PROPERTIES
TIMEOUT
45
LABELS
"integration-tests"
ENVIRONMENT
"${PRELOAD_ENV};HSA_TOOLS_LIB=$<TARGET_FILE:rocprofiler::rocprofiler-shared-library>;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}
tests/async-copy-tracing/validate.py
+61 -47
View File
@@ -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
tests/common/serialization.hpp
+16
View File
@@ -158,6 +158,13 @@ save(ArchiveT& ar, rocprofiler_hsa_api_retval_t data)
SAVE_DATA_FIELD(uint64_t_retval);
}
template <typename ArchiveT>
void
save(ArchiveT& ar, rocprofiler_marker_api_retval_t data)
{
SAVE_DATA_FIELD(int64_t_retval);
}
template <typename ArchiveT>
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 <typename ArchiveT>
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 <typename ArchiveT>
void
save(ArchiveT& ar, rocprofiler_callback_tracing_record_t data)
tests/kernel-tracing/CMakeLists.txt
+9 -9
View File
@@ -19,17 +19,17 @@ endif()
add_test(NAME test-kernel-tracing-execute COMMAND $<TARGET_FILE:reproducible-runtime>)
set(kernel-tracing-env
"${PRELOAD_ENV}"
"HSA_TOOLS_LIB=$<TARGET_FILE:rocprofiler::rocprofiler-shared-library>"
"ROCPROFILER_TOOL_OUTPUT_FILE=kernel-tracing-test.json"
"LD_LIBRARY_PATH=$<TARGET_FILE_DIR:rocprofiler::rocprofiler-shared-library>:$ENV{LD_LIBRARY_PATH}"
)
set_tests_properties(
test-kernel-tracing-execute
PROPERTIES
TIMEOUT
45
LABELS
"integration-tests"
ENVIRONMENT
"${PRELOAD_ENV};HSA_TOOLS_LIB=$<TARGET_FILE:rocprofiler::rocprofiler-shared-library>;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}
tests/kernel-tracing/validate.py
+63 -47
View File
@@ -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
tests/tools/CMakeLists.txt
+3 -10
View File
@@ -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}:$<TARGET_FILE:rocprofiler-sdk-json-tool>")
else()
set(PRELOAD_ENV "LD_PRELOAD=$<TARGET_FILE:rocprofiler-sdk-json-tool>")
endif()
tests/tools/json-tool.cpp
+99 -53
View File
@@ -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_kind_t>{
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<callback_name_info*>(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<void*>(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_kind_t>{
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<buffer_name_info*>(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<void*>(data)),
@@ -226,7 +232,8 @@ get_buffer_tracing_names()
using callback_payload_t =
std::variant<rocprofiler_callback_tracing_code_object_load_data_t,
rocprofiler_callback_tracing_code_object_kernel_symbol_register_data_t,
rocprofiler_callback_tracing_hsa_api_data_t>;
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 <typename ArchiveT>
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<code_object_callback_record_t>{};
auto kernel_symbol_records = std::deque<kernel_symbol_callback_record_t>{};
auto hsa_api_cb_records = std::deque<hsa_api_callback_record_t>{};
auto marker_api_cb_records = std::deque<marker_api_callback_record_t>{};
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<rocprofiler_callback_tracing_code_object_load_data_t*>(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<kernel_symbol_data_t*>(record.payload);
data_v.kernel_name = ::strdup(data_v.kernel_name);
auto data_v = *static_cast<kernel_symbol_data_t*>(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<rocprofiler_callback_tracing_hsa_api_data_t*>(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<rocprofiler_callback_tracing_marker_api_data_t*>(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<rocprofiler_buffer_tracing_hsa_api_record_t>{};
auto marker_api_bf_records = std::deque<rocprofiler_buffer_tracing_marker_api_record_t>{};
auto kernel_dispatch_records = std::deque<rocprofiler_buffer_tracing_kernel_dispatch_record_t>{};
auto memory_copy_records = std::deque<rocprofiler_buffer_tracing_memory_copy_record_t>{};
@@ -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<rocprofiler_buffer_tracing_marker_api_record_t*>(header->payload);
marker_api_bf_records.emplace_back(*record);
}
else if(header->kind == ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH)
{
auto* record = static_cast<rocprofiler_buffer_tracing_kernel_dispatch_record_t*>(
@@ -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<std::string_view, rocprofiler_context_id_t*>{
{"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<rocprofiler_buffer_id_t*, 3>{&api_buffered_buffer,
auto buffers = std::array<rocprofiler_buffer_id_t*, 4>{&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<std::vector<rocprofiler_agent_t>*>(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<call_stack_t*>(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<char*>(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